37#include "llvm/ADT/ArrayRef.h"
38#include "llvm/ADT/STLExtras.h"
39#include "llvm/ADT/SmallVector.h"
40#include "llvm/ADT/StringExtras.h"
41#include "llvm/ADT/StringRef.h"
42#include "llvm/ADT/Twine.h"
43#include "llvm/Frontend/HLSL/HLSLBinding.h"
44#include "llvm/Frontend/HLSL/RootSignatureValidations.h"
45#include "llvm/Support/Casting.h"
46#include "llvm/Support/DXILABI.h"
47#include "llvm/Support/ErrorHandling.h"
48#include "llvm/Support/FormatVariadic.h"
49#include "llvm/TargetParser/Triple.h"
64 case ResourceClass::SRV:
65 return RegisterType::SRV;
66 case ResourceClass::UAV:
67 return RegisterType::UAV;
68 case ResourceClass::CBuffer:
69 return RegisterType::CBuffer;
70 case ResourceClass::Sampler:
71 return RegisterType::Sampler;
73 llvm_unreachable(
"unexpected ResourceClass value");
82 case ResourceClass::SRV:
83 case ResourceClass::UAV:
85 case ResourceClass::CBuffer:
87 case ResourceClass::Sampler:
90 llvm_unreachable(
"unexpected ResourceClass value");
96 assert(RT !=
nullptr);
100 *RT = RegisterType::SRV;
104 *RT = RegisterType::UAV;
108 *RT = RegisterType::CBuffer;
112 *RT = RegisterType::Sampler;
116 *RT = RegisterType::C;
120 *RT = RegisterType::I;
129 case RegisterType::SRV:
131 case RegisterType::UAV:
133 case RegisterType::CBuffer:
135 case RegisterType::Sampler:
137 case RegisterType::C:
139 case RegisterType::I:
142 llvm_unreachable(
"unexpected RegisterType value");
147 case RegisterType::SRV:
148 return ResourceClass::SRV;
149 case RegisterType::UAV:
150 return ResourceClass::UAV;
151 case RegisterType::CBuffer:
152 return ResourceClass::CBuffer;
153 case RegisterType::Sampler:
154 return ResourceClass::Sampler;
155 case RegisterType::C:
156 case RegisterType::I:
160 llvm_unreachable(
"unexpected RegisterType value");
164 const auto *BT = dyn_cast<BuiltinType>(
Type);
168 return Builtin::BI__builtin_get_spirv_spec_constant_int;
171 switch (BT->getKind()) {
172 case BuiltinType::Bool:
173 return Builtin::BI__builtin_get_spirv_spec_constant_bool;
174 case BuiltinType::Short:
175 return Builtin::BI__builtin_get_spirv_spec_constant_short;
176 case BuiltinType::Int:
177 return Builtin::BI__builtin_get_spirv_spec_constant_int;
178 case BuiltinType::LongLong:
179 return Builtin::BI__builtin_get_spirv_spec_constant_longlong;
180 case BuiltinType::UShort:
181 return Builtin::BI__builtin_get_spirv_spec_constant_ushort;
182 case BuiltinType::UInt:
183 return Builtin::BI__builtin_get_spirv_spec_constant_uint;
184 case BuiltinType::ULongLong:
185 return Builtin::BI__builtin_get_spirv_spec_constant_ulonglong;
186 case BuiltinType::Half:
187 return Builtin::BI__builtin_get_spirv_spec_constant_half;
188 case BuiltinType::Float:
189 return Builtin::BI__builtin_get_spirv_spec_constant_float;
190 case BuiltinType::Double:
191 return Builtin::BI__builtin_get_spirv_spec_constant_double;
200 llvm::raw_svector_ostream OS(Buffer);
207 ResourceClass ResClass) {
209 "DeclBindingInfo already added");
215 DeclToBindingListIndex.try_emplace(VD, BindingsList.size());
216 return &BindingsList.emplace_back(VD, ResClass);
220 ResourceClass ResClass) {
221 auto Entry = DeclToBindingListIndex.find(VD);
222 if (Entry != DeclToBindingListIndex.end()) {
223 for (
unsigned Index = Entry->getSecond();
224 Index < BindingsList.size() && BindingsList[Index].Decl == VD;
226 if (BindingsList[Index].ResClass == ResClass)
227 return &BindingsList[Index];
234 return DeclToBindingListIndex.contains(VD);
246 getASTContext(), LexicalParent, CBuffer, KwLoc, Ident, IdentLoc, LBrace);
249 auto RC = CBuffer ? llvm::hlsl::ResourceClass::CBuffer
250 : llvm::hlsl::ResourceClass::SRV;
262 if (T->isArrayType() || T->isStructureType() || T->isConstantMatrixType())
269 assert(Context.getTypeSize(T) <= 64 &&
270 "Scalar bit widths larger than 64 not supported");
273 return Context.getTypeSize(T) / 8;
280 constexpr unsigned CBufferAlign = 16;
281 if (
const auto *RD = T->getAsRecordDecl()) {
283 for (
const FieldDecl *Field : RD->fields()) {
290 unsigned AlignSize = llvm::alignTo(Size, FieldAlign);
291 if ((AlignSize % CBufferAlign) + FieldSize > CBufferAlign) {
292 FieldAlign = CBufferAlign;
295 Size = llvm::alignTo(Size, FieldAlign);
302 unsigned ElementCount = AT->getSize().getZExtValue();
303 if (ElementCount == 0)
306 unsigned ElementSize =
308 unsigned AlignedElementSize = llvm::alignTo(ElementSize, CBufferAlign);
309 return AlignedElementSize * (ElementCount - 1) + ElementSize;
313 unsigned ElementCount = VT->getNumElements();
314 unsigned ElementSize =
316 return ElementSize * ElementCount;
319 return Context.getTypeSize(T) / 8;
330 bool HasPackOffset =
false;
331 bool HasNonPackOffset =
false;
333 VarDecl *Var = dyn_cast<VarDecl>(Field);
336 if (Field->hasAttr<HLSLPackOffsetAttr>()) {
337 PackOffsetVec.emplace_back(Var, Field->
getAttr<HLSLPackOffsetAttr>());
338 HasPackOffset =
true;
340 HasNonPackOffset =
true;
347 if (HasNonPackOffset)
354 std::sort(PackOffsetVec.begin(), PackOffsetVec.end(),
355 [](
const std::pair<VarDecl *, HLSLPackOffsetAttr *> &LHS,
356 const std::pair<VarDecl *, HLSLPackOffsetAttr *> &RHS) {
357 return LHS.second->getOffsetInBytes() <
358 RHS.second->getOffsetInBytes();
360 for (
unsigned i = 0; i < PackOffsetVec.size() - 1; i++) {
361 VarDecl *Var = PackOffsetVec[i].first;
362 HLSLPackOffsetAttr *
Attr = PackOffsetVec[i].second;
364 unsigned Begin =
Attr->getOffsetInBytes();
365 unsigned End = Begin + Size;
366 unsigned NextBegin = PackOffsetVec[i + 1].second->getOffsetInBytes();
367 if (End > NextBegin) {
368 VarDecl *NextVar = PackOffsetVec[i + 1].first;
380 CAT = dyn_cast<ConstantArrayType>(
382 return CAT !=
nullptr;
393static const HLSLAttributedResourceType *
396 "expected array of resource records");
398 while (
const ArrayType *AT = dyn_cast<ArrayType>(Ty))
400 return HLSLAttributedResourceType::findHandleTypeOnResource(Ty);
403static const HLSLAttributedResourceType *
417 return RD->isEmpty();
446 Base.getType()->castAsCXXRecordDecl()))
457 assert(RD ==
nullptr &&
458 "there should be at most 1 record by a given name in a scope");
475 Name.append(NameBaseII->
getName());
482 size_t NameLength = Name.size();
491 Name.append(llvm::Twine(suffix).str());
492 II = &AST.
Idents.
get(Name, tok::TokenKind::identifier);
499 Name.truncate(NameLength);
514 if (
const auto *CAT = dyn_cast<ConstantArrayType>(Ty)) {
516 S, CAT->getElementType()->getUnqualifiedDesugaredType());
521 CAT->getSizeModifier(),
522 CAT->getIndexTypeCVRQualifiers())
560 "struct is already HLSL buffer compatible");
574 LS->
addAttr(PackedAttr::CreateImplicit(AST));
578 if (
unsigned NumBases = StructDecl->
getNumBases()) {
579 assert(NumBases == 1 &&
"HLSL supports only one base type");
629 LS->
addAttr(PackedAttr::CreateImplicit(AST));
634 VarDecl *VD = dyn_cast<VarDecl>(D);
650 "host layout field for $Globals decl failed to be created");
667 uint32_t ImplicitBindingOrderID) {
669 HLSLResourceBindingAttr::CreateImplicit(S.
getASTContext(),
"",
"0", {});
670 Attr->setBinding(RT, std::nullopt, 0);
671 Attr->setImplicitBindingOrderID(ImplicitBindingOrderID);
678 BufDecl->setRBraceLoc(RBrace);
695 BufDecl->isCBuffer() ? RegisterType::CBuffer
705 int X,
int Y,
int Z) {
706 if (HLSLNumThreadsAttr *NT = D->
getAttr<HLSLNumThreadsAttr>()) {
707 if (NT->getX() !=
X || NT->getY() != Y || NT->getZ() != Z) {
708 Diag(NT->getLocation(), diag::err_hlsl_attribute_param_mismatch) << AL;
709 Diag(AL.
getLoc(), diag::note_conflicting_attribute);
719 int Min,
int Max,
int Preferred,
720 int SpelledArgsCount) {
721 if (HLSLWaveSizeAttr *WS = D->
getAttr<HLSLWaveSizeAttr>()) {
722 if (WS->getMin() !=
Min || WS->getMax() !=
Max ||
723 WS->getPreferred() != Preferred ||
724 WS->getSpelledArgsCount() != SpelledArgsCount) {
725 Diag(WS->getLocation(), diag::err_hlsl_attribute_param_mismatch) << AL;
726 Diag(AL.
getLoc(), diag::note_conflicting_attribute);
732 Result->setSpelledArgsCount(SpelledArgsCount);
736HLSLVkConstantIdAttr *
742 Diag(AL.
getLoc(), diag::warn_attribute_ignored) << AL;
750 Diag(VD->getLocation(), diag::err_specialization_const);
754 if (!VD->getType().isConstQualified()) {
755 Diag(VD->getLocation(), diag::err_specialization_const);
759 if (HLSLVkConstantIdAttr *CI = D->
getAttr<HLSLVkConstantIdAttr>()) {
760 if (CI->getId() != Id) {
761 Diag(CI->getLocation(), diag::err_hlsl_attribute_param_mismatch) << AL;
762 Diag(AL.
getLoc(), diag::note_conflicting_attribute);
767 HLSLVkConstantIdAttr *
Result =
774 llvm::Triple::EnvironmentType ShaderType) {
775 if (HLSLShaderAttr *NT = D->
getAttr<HLSLShaderAttr>()) {
776 if (NT->getType() != ShaderType) {
777 Diag(NT->getLocation(), diag::err_hlsl_attribute_param_mismatch) << AL;
778 Diag(AL.
getLoc(), diag::note_conflicting_attribute);
782 return HLSLShaderAttr::Create(
getASTContext(), ShaderType, AL);
785HLSLParamModifierAttr *
787 HLSLParamModifierAttr::Spelling Spelling) {
790 if (HLSLParamModifierAttr *PA = D->
getAttr<HLSLParamModifierAttr>()) {
791 if ((PA->isIn() && Spelling == HLSLParamModifierAttr::Keyword_out) ||
792 (PA->isOut() && Spelling == HLSLParamModifierAttr::Keyword_in)) {
793 D->
dropAttr<HLSLParamModifierAttr>();
795 return HLSLParamModifierAttr::Create(
797 HLSLParamModifierAttr::Keyword_inout);
799 Diag(AL.
getLoc(), diag::err_hlsl_duplicate_parameter_modifier) << AL;
800 Diag(PA->getLocation(), diag::note_conflicting_attribute);
826 if (HLSLShaderAttr::isValidShaderType(Env) && Env != llvm::Triple::Library) {
827 if (
const auto *Shader = FD->
getAttr<HLSLShaderAttr>()) {
830 if (Shader->getType() != Env) {
831 Diag(Shader->getLocation(), diag::err_hlsl_entry_shader_attr_mismatch)
843 case llvm::Triple::UnknownEnvironment:
844 case llvm::Triple::Library:
846 case llvm::Triple::RootSignature:
847 llvm_unreachable(
"rootsig environment has no functions");
849 llvm_unreachable(
"Unhandled environment in triple");
855 HLSLAppliedSemanticAttr *Semantic,
860 const auto *ShaderAttr = FD->
getAttr<HLSLShaderAttr>();
861 assert(ShaderAttr &&
"Entry point has no shader attribute");
862 llvm::Triple::EnvironmentType ST = ShaderAttr->getType();
863 auto SemanticName = Semantic->getSemanticName().upper();
868 if (SemanticName ==
"SV_POSITION") {
869 return (ST == llvm::Triple::Vertex && !IsInput) ||
870 (ST == llvm::Triple::Pixel && IsInput);
872 if (SemanticName ==
"SV_VERTEXID")
878bool SemaHLSL::determineActiveSemanticOnScalar(
FunctionDecl *FD,
881 SemanticInfo &ActiveSemantic,
882 SemaHLSL::SemanticContext &SC) {
883 if (ActiveSemantic.Semantic ==
nullptr) {
884 ActiveSemantic.Semantic = D->
getAttr<HLSLParsedSemanticAttr>();
885 if (ActiveSemantic.Semantic)
886 ActiveSemantic.Index = ActiveSemantic.Semantic->getSemanticIndex();
889 if (!ActiveSemantic.Semantic) {
895 HLSLAppliedSemanticAttr(
getASTContext(), *ActiveSemantic.Semantic,
896 ActiveSemantic.Semantic->getAttrName()->getName(),
897 ActiveSemantic.Index.value_or(0));
901 checkSemanticAnnotation(FD, D, A, SC);
902 OutputDecl->addAttr(A);
904 unsigned Location = ActiveSemantic.Index.value_or(0);
907 SC.CurrentIOType & IOType::In)) {
908 bool HasVkLocation =
false;
909 if (
auto *A = D->getAttr<HLSLVkLocationAttr>()) {
910 HasVkLocation = true;
911 Location = A->getLocation();
914 if (SC.UsesExplicitVkLocations.value_or(HasVkLocation) != HasVkLocation) {
915 Diag(D->getLocation(), diag::err_hlsl_semantic_partial_explicit_indexing);
918 SC.UsesExplicitVkLocations = HasVkLocation;
921 const ConstantArrayType *AT = dyn_cast<ConstantArrayType>(D->getType());
922 unsigned ElementCount = AT ? AT->
getZExtSize() : 1;
923 ActiveSemantic.Index = Location + ElementCount;
925 Twine BaseName = Twine(ActiveSemantic.Semantic->getAttrName()->getName());
926 for (
unsigned I = 0; I < ElementCount; ++I) {
927 Twine VariableName = BaseName.concat(Twine(Location + I));
929 auto [_, Inserted] = SC.ActiveSemantics.insert(VariableName.str());
931 Diag(D->getLocation(), diag::err_hlsl_semantic_index_overlap)
932 << VariableName.str();
943 SemanticInfo &ActiveSemantic,
944 SemaHLSL::SemanticContext &SC) {
945 if (ActiveSemantic.Semantic ==
nullptr) {
946 ActiveSemantic.Semantic = D->
getAttr<HLSLParsedSemanticAttr>();
947 if (ActiveSemantic.Semantic)
948 ActiveSemantic.Index = ActiveSemantic.Semantic->getSemanticIndex();
954 const RecordType *RT = dyn_cast<RecordType>(T);
956 return determineActiveSemanticOnScalar(FD, OutputDecl, D, ActiveSemantic,
959 const RecordDecl *RD = RT->getDecl();
960 for (FieldDecl *Field : RD->
fields()) {
961 SemanticInfo Info = ActiveSemantic;
962 if (!determineActiveSemantic(FD, OutputDecl, Field, Info, SC)) {
963 Diag(
Field->getLocation(), diag::note_hlsl_semantic_used_here) <<
Field;
966 if (ActiveSemantic.Semantic)
967 ActiveSemantic = Info;
974 const auto *ShaderAttr = FD->
getAttr<HLSLShaderAttr>();
975 assert(ShaderAttr &&
"Entry point has no shader attribute");
976 llvm::Triple::EnvironmentType ST = ShaderAttr->getType();
980 case llvm::Triple::Pixel:
981 case llvm::Triple::Vertex:
982 case llvm::Triple::Geometry:
983 case llvm::Triple::Hull:
984 case llvm::Triple::Domain:
985 case llvm::Triple::RayGeneration:
986 case llvm::Triple::Intersection:
987 case llvm::Triple::AnyHit:
988 case llvm::Triple::ClosestHit:
989 case llvm::Triple::Miss:
990 case llvm::Triple::Callable:
991 if (
const auto *NT = FD->
getAttr<HLSLNumThreadsAttr>()) {
992 diagnoseAttrStageMismatch(NT, ST,
993 {llvm::Triple::Compute,
994 llvm::Triple::Amplification,
995 llvm::Triple::Mesh});
998 if (
const auto *WS = FD->
getAttr<HLSLWaveSizeAttr>()) {
999 diagnoseAttrStageMismatch(WS, ST,
1000 {llvm::Triple::Compute,
1001 llvm::Triple::Amplification,
1002 llvm::Triple::Mesh});
1007 case llvm::Triple::Compute:
1008 case llvm::Triple::Amplification:
1009 case llvm::Triple::Mesh:
1010 if (!FD->
hasAttr<HLSLNumThreadsAttr>()) {
1012 << llvm::Triple::getEnvironmentTypeName(ST);
1015 if (
const auto *WS = FD->
getAttr<HLSLWaveSizeAttr>()) {
1016 if (Ver < VersionTuple(6, 6)) {
1017 Diag(WS->getLocation(), diag::err_hlsl_attribute_in_wrong_shader_model)
1020 }
else if (WS->getSpelledArgsCount() > 1 && Ver < VersionTuple(6, 8)) {
1023 diag::err_hlsl_attribute_number_arguments_insufficient_shader_model)
1024 << WS << WS->getSpelledArgsCount() <<
"6.8";
1029 case llvm::Triple::RootSignature:
1030 llvm_unreachable(
"rootsig environment has no function entry point");
1032 llvm_unreachable(
"Unhandled environment in triple");
1035 SemaHLSL::SemanticContext InputSC = {};
1036 InputSC.CurrentIOType = IOType::In;
1039 SemanticInfo ActiveSemantic;
1040 ActiveSemantic.Semantic = Param->getAttr<HLSLParsedSemanticAttr>();
1041 if (ActiveSemantic.Semantic)
1042 ActiveSemantic.Index = ActiveSemantic.Semantic->getSemanticIndex();
1045 if (!determineActiveSemantic(FD, Param, Param, ActiveSemantic, InputSC)) {
1046 Diag(Param->getLocation(), diag::note_previous_decl) << Param;
1051 SemanticInfo ActiveSemantic;
1052 SemaHLSL::SemanticContext OutputSC = {};
1053 OutputSC.CurrentIOType = IOType::Out;
1054 ActiveSemantic.Semantic = FD->
getAttr<HLSLParsedSemanticAttr>();
1055 if (ActiveSemantic.Semantic)
1056 ActiveSemantic.Index = ActiveSemantic.Semantic->getSemanticIndex();
1058 determineActiveSemantic(FD, FD, FD, ActiveSemantic, OutputSC);
1061void SemaHLSL::checkSemanticAnnotation(
1063 const HLSLAppliedSemanticAttr *SemanticAttr,
const SemanticContext &SC) {
1064 auto *ShaderAttr = EntryPoint->
getAttr<HLSLShaderAttr>();
1065 assert(ShaderAttr &&
"Entry point has no shader attribute");
1066 llvm::Triple::EnvironmentType ST = ShaderAttr->getType();
1068 auto SemanticName = SemanticAttr->getSemanticName().upper();
1069 if (SemanticName ==
"SV_DISPATCHTHREADID" ||
1070 SemanticName ==
"SV_GROUPINDEX" || SemanticName ==
"SV_GROUPTHREADID" ||
1071 SemanticName ==
"SV_GROUPID") {
1073 if (ST != llvm::Triple::Compute)
1074 diagnoseSemanticStageMismatch(SemanticAttr, ST, SC.CurrentIOType,
1075 {{llvm::Triple::Compute, IOType::In}});
1077 if (SemanticAttr->getSemanticIndex() != 0) {
1078 std::string PrettyName =
1079 "'" + SemanticAttr->getSemanticName().str() +
"'";
1080 Diag(SemanticAttr->getLoc(),
1081 diag::err_hlsl_semantic_indexing_not_supported)
1087 if (SemanticName ==
"SV_POSITION") {
1090 diagnoseSemanticStageMismatch(SemanticAttr, ST, SC.CurrentIOType,
1091 {{llvm::Triple::Vertex, IOType::InOut},
1092 {llvm::Triple::Pixel, IOType::In}});
1095 if (SemanticName ==
"SV_VERTEXID") {
1096 diagnoseSemanticStageMismatch(SemanticAttr, ST, SC.CurrentIOType,
1097 {{llvm::Triple::Vertex, IOType::In}});
1101 if (SemanticName ==
"SV_TARGET") {
1102 diagnoseSemanticStageMismatch(SemanticAttr, ST, SC.CurrentIOType,
1103 {{llvm::Triple::Pixel, IOType::Out}});
1109 if (SemanticAttr->getAttrName()->getName().starts_with_insensitive(
"SV_"))
1110 llvm_unreachable(
"Unknown SemanticAttr");
1113void SemaHLSL::diagnoseAttrStageMismatch(
1114 const Attr *A, llvm::Triple::EnvironmentType Stage,
1115 std::initializer_list<llvm::Triple::EnvironmentType> AllowedStages) {
1116 SmallVector<StringRef, 8> StageStrings;
1117 llvm::transform(AllowedStages, std::back_inserter(StageStrings),
1118 [](llvm::Triple::EnvironmentType ST) {
1120 HLSLShaderAttr::ConvertEnvironmentTypeToStr(ST));
1122 Diag(A->
getLoc(), diag::err_hlsl_attr_unsupported_in_stage)
1123 << A->
getAttrName() << llvm::Triple::getEnvironmentTypeName(Stage)
1124 << (AllowedStages.size() != 1) << join(StageStrings,
", ");
1127void SemaHLSL::diagnoseSemanticStageMismatch(
1128 const Attr *A, llvm::Triple::EnvironmentType Stage, IOType CurrentIOType,
1129 std::initializer_list<SemanticStageInfo> Allowed) {
1131 for (
auto &Case : Allowed) {
1132 if (Case.Stage != Stage)
1135 if (CurrentIOType & Case.AllowedIOTypesMask)
1138 SmallVector<std::string, 8> ValidCases;
1140 Allowed, std::back_inserter(ValidCases), [](SemanticStageInfo Case) {
1141 SmallVector<std::string, 2> ValidType;
1142 if (Case.AllowedIOTypesMask & IOType::In)
1143 ValidType.push_back(
"input");
1144 if (Case.AllowedIOTypesMask & IOType::Out)
1145 ValidType.push_back(
"output");
1147 HLSLShaderAttr::ConvertEnvironmentTypeToStr(Case.Stage)) +
1148 " " + join(ValidType,
"/");
1150 Diag(A->
getLoc(), diag::err_hlsl_semantic_unsupported_iotype_for_stage)
1151 << A->
getAttrName() << (CurrentIOType & IOType::In ?
"input" :
"output")
1152 << llvm::Triple::getEnvironmentTypeName(Case.Stage)
1153 << join(ValidCases,
", ");
1157 SmallVector<StringRef, 8> StageStrings;
1159 Allowed, std::back_inserter(StageStrings), [](SemanticStageInfo Case) {
1161 HLSLShaderAttr::ConvertEnvironmentTypeToStr(Case.Stage));
1164 Diag(A->
getLoc(), diag::err_hlsl_attr_unsupported_in_stage)
1165 << A->
getAttrName() << llvm::Triple::getEnvironmentTypeName(Stage)
1166 << (Allowed.size() != 1) << join(StageStrings,
", ");
1169template <CastKind Kind>
1172 Ty = VTy->getElementType();
1177template <CastKind Kind>
1189 if (LHSFloat && RHSFloat) {
1217 if (LHSSigned == RHSSigned) {
1218 if (IsCompAssign || IntOrder >= 0)
1226 if (IntOrder != (LHSSigned ? 1 : -1)) {
1227 if (IsCompAssign || RHSSigned)
1235 if (Ctx.getIntWidth(LElTy) != Ctx.getIntWidth(RElTy)) {
1236 if (IsCompAssign || LHSSigned)
1252 QualType ElTy = Ctx.getCorrespondingUnsignedType(LHSSigned ? LElTy : RElTy);
1253 QualType NewTy = Ctx.getExtVectorType(
1263 return CK_FloatingCast;
1265 return CK_IntegralCast;
1267 return CK_IntegralToFloating;
1269 return CK_FloatingToIntegral;
1275 bool IsCompAssign) {
1282 if (!LVecTy && IsCompAssign) {
1284 RHS =
SemaRef.ImpCastExprToType(RHS.
get(), RElTy, CK_HLSLVectorTruncation);
1286 if (Ctx.hasSameUnqualifiedType(LHSType, RHSType))
1288 RHS =
SemaRef.ImpCastExprToType(RHS.
get(), LHSType,
1293 unsigned EndSz = std::numeric_limits<unsigned>::max();
1296 LSz = EndSz = LVecTy->getNumElements();
1299 assert(EndSz != std::numeric_limits<unsigned>::max() &&
1300 "one of the above should have had a value");
1304 if (IsCompAssign && LSz != EndSz) {
1306 diag::err_hlsl_vector_compound_assignment_truncation)
1307 << LHSType << RHSType;
1313 if (!IsCompAssign && LVecTy && LVecTy->getNumElements() > EndSz)
1318 if (!IsCompAssign && !LVecTy)
1322 if (Ctx.hasSameUnqualifiedType(LHSType, RHSType))
1323 return Ctx.getCommonSugaredType(LHSType, RHSType);
1331 LElTy, RElTy, IsCompAssign);
1334 "HLSL Vectors can only contain integer or floating point types");
1336 LElTy, RElTy, IsCompAssign);
1341 assert((Opc == BO_LOr || Opc == BO_LAnd) &&
1342 "Called with non-logical operator");
1344 llvm::raw_svector_ostream OS(Buff);
1346 StringRef NewFnName = Opc == BO_LOr ?
"or" :
"and";
1347 OS << NewFnName <<
"(";
1357std::pair<IdentifierInfo *, bool>
1360 std::string IdStr =
"__hlsl_rootsig_decl_" + std::to_string(Hash);
1367 return {DeclIdent,
Found};
1378 for (
auto &RootSigElement : RootElements)
1379 Elements.push_back(RootSigElement.getElement());
1383 DeclIdent,
SemaRef.getLangOpts().HLSLRootSigVer, Elements);
1385 SignatureDecl->setImplicit();
1391 if (RootSigOverrideIdent) {
1394 if (
SemaRef.LookupQualifiedName(R, DC))
1395 return dyn_cast<HLSLRootSignatureDecl>(R.getFoundDecl());
1403struct PerVisibilityBindingChecker {
1406 std::array<llvm::hlsl::BindingInfoBuilder, 8> Builders;
1410 llvm::dxbc::ShaderVisibility Vis;
1415 PerVisibilityBindingChecker(
SemaHLSL *S) : S(S) {}
1417 void trackBinding(llvm::dxbc::ShaderVisibility
Visibility,
1418 llvm::dxil::ResourceClass RC, uint32_t Space,
1419 uint32_t LowerBound, uint32_t UpperBound,
1420 const hlsl::RootSignatureElement *Elem) {
1422 assert(BuilderIndex < Builders.size() &&
1423 "Not enough builders for visibility type");
1424 Builders[BuilderIndex].trackBinding(RC, Space, LowerBound, UpperBound,
1425 static_cast<const void *
>(Elem));
1427 static_assert(llvm::to_underlying(llvm::dxbc::ShaderVisibility::All) == 0,
1428 "'All' visibility must come first");
1429 if (
Visibility == llvm::dxbc::ShaderVisibility::All)
1430 for (
size_t I = 1, E = Builders.size(); I < E; ++I)
1431 Builders[I].trackBinding(RC, Space, LowerBound, UpperBound,
1432 static_cast<const void *
>(Elem));
1434 ElemInfoMap.push_back({Elem,
Visibility,
false});
1437 ElemInfo &
getInfo(
const hlsl::RootSignatureElement *Elem) {
1438 auto It = llvm::lower_bound(
1440 [](
const auto &LHS,
const auto &RHS) {
return LHS.Elem < RHS; });
1441 assert(It->Elem == Elem &&
"Element not in map");
1445 bool checkOverlap() {
1446 llvm::sort(ElemInfoMap, [](
const auto &LHS,
const auto &RHS) {
1447 return LHS.Elem < RHS.Elem;
1450 bool HadOverlap =
false;
1452 using llvm::hlsl::BindingInfoBuilder;
1453 auto ReportOverlap = [
this,
1454 &HadOverlap](
const BindingInfoBuilder &Builder,
1455 const llvm::hlsl::Binding &Reported) {
1459 static_cast<const hlsl::RootSignatureElement *
>(Reported.Cookie);
1460 const llvm::hlsl::Binding &
Previous = Builder.findOverlapping(Reported);
1461 const auto *PrevElem =
1462 static_cast<const hlsl::RootSignatureElement *
>(
Previous.Cookie);
1464 ElemInfo &Info =
getInfo(Elem);
1469 Info.Diagnosed =
true;
1471 ElemInfo &PrevInfo =
getInfo(PrevElem);
1472 llvm::dxbc::ShaderVisibility CommonVis =
1473 Info.Vis == llvm::dxbc::ShaderVisibility::All ? PrevInfo.Vis
1476 this->S->
Diag(Elem->
getLocation(), diag::err_hlsl_resource_range_overlap)
1477 << llvm::to_underlying(Reported.RC) << Reported.LowerBound
1478 << Reported.isUnbounded() << Reported.UpperBound
1483 this->S->
Diag(PrevElem->getLocation(),
1484 diag::note_hlsl_resource_range_here);
1487 for (BindingInfoBuilder &Builder : Builders)
1488 Builder.calculateBindingInfo(ReportOverlap);
1508 bool HadError =
false;
1509 auto ReportError = [
this, &HadError](
SourceLocation Loc, uint32_t LowerBound,
1510 uint32_t UpperBound) {
1512 this->
Diag(Loc, diag::err_hlsl_invalid_rootsig_value)
1513 << LowerBound << UpperBound;
1520 this->
Diag(Loc, diag::err_hlsl_invalid_rootsig_value)
1521 << llvm::formatv(
"{0:f}", LowerBound).sstr<6>()
1522 << llvm::formatv(
"{0:f}", UpperBound).sstr<6>();
1525 auto VerifyRegister = [ReportError](
SourceLocation Loc, uint32_t Register) {
1526 if (!llvm::hlsl::rootsig::verifyRegisterValue(Register))
1527 ReportError(Loc, 0, 0xfffffffe);
1530 auto VerifySpace = [ReportError](
SourceLocation Loc, uint32_t Space) {
1531 if (!llvm::hlsl::rootsig::verifyRegisterSpace(Space))
1532 ReportError(Loc, 0, 0xffffffef);
1535 const uint32_t Version =
1536 llvm::to_underlying(
SemaRef.getLangOpts().HLSLRootSigVer);
1537 const uint32_t VersionEnum = Version - 1;
1538 auto ReportFlagError = [
this, &HadError, VersionEnum](
SourceLocation Loc) {
1540 this->
Diag(Loc, diag::err_hlsl_invalid_rootsig_flag)
1547 const llvm::hlsl::rootsig::RootElement &Elem = RootSigElem.
getElement();
1548 if (
const auto *Descriptor =
1549 std::get_if<llvm::hlsl::rootsig::RootDescriptor>(&Elem)) {
1550 VerifyRegister(Loc, Descriptor->Reg.Number);
1551 VerifySpace(Loc, Descriptor->Space);
1553 if (!llvm::hlsl::rootsig::verifyRootDescriptorFlag(Version,
1555 ReportFlagError(Loc);
1556 }
else if (
const auto *Constants =
1557 std::get_if<llvm::hlsl::rootsig::RootConstants>(&Elem)) {
1558 VerifyRegister(Loc, Constants->Reg.Number);
1559 VerifySpace(Loc, Constants->Space);
1560 }
else if (
const auto *Sampler =
1561 std::get_if<llvm::hlsl::rootsig::StaticSampler>(&Elem)) {
1562 VerifyRegister(Loc, Sampler->Reg.Number);
1563 VerifySpace(Loc, Sampler->Space);
1566 "By construction, parseFloatParam can't produce a NaN from a "
1567 "float_literal token");
1569 if (!llvm::hlsl::rootsig::verifyMaxAnisotropy(Sampler->MaxAnisotropy))
1570 ReportError(Loc, 0, 16);
1571 if (!llvm::hlsl::rootsig::verifyMipLODBias(Sampler->MipLODBias))
1572 ReportFloatError(Loc, -16.f, 15.99f);
1573 }
else if (
const auto *Clause =
1574 std::get_if<llvm::hlsl::rootsig::DescriptorTableClause>(
1576 VerifyRegister(Loc, Clause->Reg.Number);
1577 VerifySpace(Loc, Clause->Space);
1579 if (!llvm::hlsl::rootsig::verifyNumDescriptors(Clause->NumDescriptors)) {
1583 ReportError(Loc, 1, 0xfffffffe);
1586 if (!llvm::hlsl::rootsig::verifyDescriptorRangeFlag(Version, Clause->Type,
1588 ReportFlagError(Loc);
1592 PerVisibilityBindingChecker BindingChecker(
this);
1593 SmallVector<std::pair<
const llvm::hlsl::rootsig::DescriptorTableClause *,
1598 const llvm::hlsl::rootsig::RootElement &Elem = RootSigElem.
getElement();
1599 if (
const auto *Descriptor =
1600 std::get_if<llvm::hlsl::rootsig::RootDescriptor>(&Elem)) {
1601 uint32_t LowerBound(Descriptor->Reg.Number);
1602 uint32_t UpperBound(LowerBound);
1604 BindingChecker.trackBinding(
1605 Descriptor->Visibility,
1606 static_cast<llvm::dxil::ResourceClass
>(Descriptor->Type),
1607 Descriptor->Space, LowerBound, UpperBound, &RootSigElem);
1608 }
else if (
const auto *Constants =
1609 std::get_if<llvm::hlsl::rootsig::RootConstants>(&Elem)) {
1610 uint32_t LowerBound(Constants->Reg.Number);
1611 uint32_t UpperBound(LowerBound);
1613 BindingChecker.trackBinding(
1614 Constants->Visibility, llvm::dxil::ResourceClass::CBuffer,
1615 Constants->Space, LowerBound, UpperBound, &RootSigElem);
1616 }
else if (
const auto *Sampler =
1617 std::get_if<llvm::hlsl::rootsig::StaticSampler>(&Elem)) {
1618 uint32_t LowerBound(Sampler->Reg.Number);
1619 uint32_t UpperBound(LowerBound);
1621 BindingChecker.trackBinding(
1622 Sampler->Visibility, llvm::dxil::ResourceClass::Sampler,
1623 Sampler->Space, LowerBound, UpperBound, &RootSigElem);
1624 }
else if (
const auto *Clause =
1625 std::get_if<llvm::hlsl::rootsig::DescriptorTableClause>(
1628 UnboundClauses.emplace_back(Clause, &RootSigElem);
1629 }
else if (
const auto *Table =
1630 std::get_if<llvm::hlsl::rootsig::DescriptorTable>(&Elem)) {
1631 assert(UnboundClauses.size() == Table->NumClauses &&
1632 "Number of unbound elements must match the number of clauses");
1633 bool HasAnySampler =
false;
1634 bool HasAnyNonSampler =
false;
1635 uint64_t Offset = 0;
1636 bool IsPrevUnbound =
false;
1637 for (
const auto &[Clause, ClauseElem] : UnboundClauses) {
1639 if (Clause->Type == llvm::dxil::ResourceClass::Sampler)
1640 HasAnySampler =
true;
1642 HasAnyNonSampler =
true;
1644 if (HasAnySampler && HasAnyNonSampler)
1645 Diag(Loc, diag::err_hlsl_invalid_mixed_resources);
1650 if (Clause->NumDescriptors == 0)
1654 Clause->Offset == llvm::hlsl::rootsig::DescriptorTableOffsetAppend;
1656 Offset = Clause->Offset;
1658 uint64_t RangeBound = llvm::hlsl::rootsig::computeRangeBound(
1659 Offset, Clause->NumDescriptors);
1661 if (IsPrevUnbound && IsAppending)
1662 Diag(Loc, diag::err_hlsl_appending_onto_unbound);
1663 else if (!llvm::hlsl::rootsig::verifyNoOverflowedOffset(RangeBound))
1664 Diag(Loc, diag::err_hlsl_offset_overflow) << Offset << RangeBound;
1667 Offset = RangeBound + 1;
1668 IsPrevUnbound = Clause->NumDescriptors ==
1669 llvm::hlsl::rootsig::NumDescriptorsUnbounded;
1672 uint32_t LowerBound(Clause->Reg.Number);
1673 uint32_t UpperBound = llvm::hlsl::rootsig::computeRangeBound(
1674 LowerBound, Clause->NumDescriptors);
1676 BindingChecker.trackBinding(
1678 static_cast<llvm::dxil::ResourceClass
>(Clause->Type), Clause->Space,
1679 LowerBound, UpperBound, ClauseElem);
1681 UnboundClauses.clear();
1685 return BindingChecker.checkOverlap();
1690 Diag(AL.
getLoc(), diag::err_attribute_wrong_number_arguments) << AL << 1;
1695 if (
auto *RS = D->
getAttr<RootSignatureAttr>()) {
1696 if (RS->getSignatureIdent() != Ident) {
1697 Diag(AL.
getLoc(), diag::err_disallowed_duplicate_attribute) << RS;
1701 Diag(AL.
getLoc(), diag::warn_duplicate_attribute_exact) << RS;
1707 if (
auto *SignatureDecl =
1708 dyn_cast<HLSLRootSignatureDecl>(R.getFoundDecl())) {
1715 llvm::VersionTuple SMVersion =
1720 uint32_t ZMax = 1024;
1721 uint32_t ThreadMax = 1024;
1722 if (IsDXIL && SMVersion.getMajor() <= 4) {
1725 }
else if (IsDXIL && SMVersion.getMajor() == 5) {
1735 diag::err_hlsl_numthreads_argument_oor)
1744 diag::err_hlsl_numthreads_argument_oor)
1753 diag::err_hlsl_numthreads_argument_oor)
1758 if (
X * Y * Z > ThreadMax) {
1759 Diag(AL.
getLoc(), diag::err_hlsl_numthreads_invalid) << ThreadMax;
1776 if (SpelledArgsCount == 0 || SpelledArgsCount > 3)
1784 if (SpelledArgsCount > 1 &&
1788 uint32_t Preferred = 0;
1789 if (SpelledArgsCount > 2 &&
1793 if (SpelledArgsCount > 2) {
1796 diag::err_attribute_power_of_two_in_range)
1797 << AL << llvm::dxil::MinWaveSize << llvm::dxil::MaxWaveSize
1802 if (Preferred < Min || Preferred >
Max) {
1804 diag::err_attribute_power_of_two_in_range)
1805 << AL <<
Min <<
Max << Preferred;
1808 }
else if (SpelledArgsCount > 1) {
1811 diag::err_attribute_power_of_two_in_range)
1812 << AL << llvm::dxil::MinWaveSize << llvm::dxil::MaxWaveSize <<
Max;
1816 Diag(AL.
getLoc(), diag::err_attribute_argument_invalid) << AL << 1;
1819 Diag(AL.
getLoc(), diag::warn_attr_min_eq_max) << AL;
1824 diag::err_attribute_power_of_two_in_range)
1825 << AL << llvm::dxil::MinWaveSize << llvm::dxil::MaxWaveSize <<
Min;
1830 HLSLWaveSizeAttr *NewAttr =
1867 uint32_t Binding = 0;
1891 if (!T->hasUnsignedIntegerRepresentation() ||
1892 (VT && VT->getNumElements() > 3)) {
1893 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type)
1894 << AL <<
"uint/uint2/uint3";
1903 if (!T->hasFloatingRepresentation() || (VT && VT->getNumElements() > 4)) {
1904 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type)
1905 << AL <<
"float/float1/float2/float3/float4";
1913 std::optional<unsigned> Index) {
1917 QualType ValueType = VD->getType();
1918 if (
auto *FD = dyn_cast<FunctionDecl>(D))
1921 bool IsOutput =
false;
1922 if (HLSLParamModifierAttr *MA = D->
getAttr<HLSLParamModifierAttr>()) {
1929 if (SemanticName ==
"SV_DISPATCHTHREADID") {
1932 Diag(AL.
getLoc(), diag::err_hlsl_semantic_output_not_supported) << AL;
1933 if (Index.has_value())
1934 Diag(AL.
getLoc(), diag::err_hlsl_semantic_indexing_not_supported) << AL;
1939 if (SemanticName ==
"SV_GROUPINDEX") {
1941 Diag(AL.
getLoc(), diag::err_hlsl_semantic_output_not_supported) << AL;
1942 if (Index.has_value())
1943 Diag(AL.
getLoc(), diag::err_hlsl_semantic_indexing_not_supported) << AL;
1948 if (SemanticName ==
"SV_GROUPTHREADID") {
1951 Diag(AL.
getLoc(), diag::err_hlsl_semantic_output_not_supported) << AL;
1952 if (Index.has_value())
1953 Diag(AL.
getLoc(), diag::err_hlsl_semantic_indexing_not_supported) << AL;
1958 if (SemanticName ==
"SV_GROUPID") {
1961 Diag(AL.
getLoc(), diag::err_hlsl_semantic_output_not_supported) << AL;
1962 if (Index.has_value())
1963 Diag(AL.
getLoc(), diag::err_hlsl_semantic_indexing_not_supported) << AL;
1968 if (SemanticName ==
"SV_POSITION") {
1969 const auto *VT = ValueType->getAs<
VectorType>();
1970 if (!ValueType->hasFloatingRepresentation() ||
1971 (VT && VT->getNumElements() > 4))
1972 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type)
1973 << AL <<
"float/float1/float2/float3/float4";
1978 if (SemanticName ==
"SV_VERTEXID") {
1979 uint64_t SizeInBits =
SemaRef.Context.getTypeSize(ValueType);
1980 if (!ValueType->isUnsignedIntegerType() || SizeInBits != 32)
1981 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type) << AL <<
"uint";
1986 if (SemanticName ==
"SV_TARGET") {
1987 const auto *VT = ValueType->getAs<
VectorType>();
1988 if (!ValueType->hasFloatingRepresentation() ||
1989 (VT && VT->getNumElements() > 4))
1990 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type)
1991 << AL <<
"float/float1/float2/float3/float4";
1996 Diag(AL.
getLoc(), diag::err_hlsl_unknown_semantic) << AL;
2000 uint32_t IndexValue(0), ExplicitIndex(0);
2003 assert(0 &&
"HLSLUnparsedSemantic is expected to have 2 int arguments.");
2005 assert(IndexValue > 0 ? ExplicitIndex :
true);
2006 std::optional<unsigned> Index =
2007 ExplicitIndex ? std::optional<unsigned>(IndexValue) : std::nullopt;
2017 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_ast_node)
2018 << AL <<
"shader constant in a constant buffer";
2022 uint32_t SubComponent;
2032 bool IsAggregateTy = (T->isArrayType() || T->isStructureType());
2037 if (IsAggregateTy) {
2038 Diag(AL.
getLoc(), diag::err_hlsl_invalid_register_or_packoffset);
2042 if ((Component * 32 + Size) > 128) {
2043 Diag(AL.
getLoc(), diag::err_hlsl_packoffset_cross_reg_boundary);
2048 EltTy = VT->getElementType();
2050 if (Align > 32 && Component == 1) {
2053 Diag(AL.
getLoc(), diag::err_hlsl_packoffset_alignment_mismatch)
2067 if (!
SemaRef.checkStringLiteralArgumentAttr(AL, 0, Str, &ArgLoc))
2070 llvm::Triple::EnvironmentType ShaderType;
2071 if (!HLSLShaderAttr::ConvertStrToEnvironmentType(Str, ShaderType)) {
2072 Diag(AL.
getLoc(), diag::warn_attribute_type_not_supported)
2073 << AL << Str << ArgLoc;
2087 assert(AttrList.size() &&
"expected list of resource attributes");
2094 HLSLAttributedResourceType::Attributes ResAttrs;
2096 bool HasResourceClass =
false;
2097 bool HasResourceDimension =
false;
2098 for (
const Attr *A : AttrList) {
2103 case attr::HLSLResourceClass: {
2105 if (HasResourceClass) {
2107 ? diag::warn_duplicate_attribute_exact
2108 : diag::warn_duplicate_attribute)
2112 ResAttrs.ResourceClass = RC;
2113 HasResourceClass =
true;
2116 case attr::HLSLResourceDimension: {
2117 llvm::dxil::ResourceDimension RD =
2119 if (HasResourceDimension) {
2121 ? diag::warn_duplicate_attribute_exact
2122 : diag::warn_duplicate_attribute)
2126 ResAttrs.ResourceDimension = RD;
2127 HasResourceDimension =
true;
2131 if (ResAttrs.IsROV) {
2135 ResAttrs.IsROV =
true;
2137 case attr::HLSLRawBuffer:
2138 if (ResAttrs.RawBuffer) {
2142 ResAttrs.RawBuffer =
true;
2144 case attr::HLSLIsCounter:
2145 if (ResAttrs.IsCounter) {
2149 ResAttrs.IsCounter =
true;
2151 case attr::HLSLContainedType: {
2154 if (!ContainedTy.
isNull()) {
2156 ? diag::warn_duplicate_attribute_exact
2157 : diag::warn_duplicate_attribute)
2166 llvm_unreachable(
"unhandled resource attribute type");
2170 if (!HasResourceClass) {
2171 S.
Diag(AttrList.back()->getRange().getEnd(),
2172 diag::err_hlsl_missing_resource_class);
2177 Wrapped, ContainedTy, ResAttrs);
2179 if (LocInfo && ContainedTyInfo) {
2192 if (!T->isHLSLResourceType()) {
2193 Diag(AL.
getLoc(), diag::err_hlsl_attribute_needs_intangible_type)
2208 AttributeCommonInfo::AS_CXX11, 0, false ,
2213 case ParsedAttr::AT_HLSLResourceClass: {
2215 Diag(AL.
getLoc(), diag::err_attribute_argument_type)
2226 if (!HLSLResourceClassAttr::ConvertStrToResourceClass(Identifier, RC)) {
2227 Diag(ArgLoc, diag::warn_attribute_type_not_supported)
2228 <<
"ResourceClass" << Identifier;
2231 A = HLSLResourceClassAttr::Create(
getASTContext(), RC, ACI);
2235 case ParsedAttr::AT_HLSLResourceDimension: {
2236 StringRef Identifier;
2238 if (!
SemaRef.checkStringLiteralArgumentAttr(AL, 0, Identifier, &ArgLoc))
2242 llvm::dxil::ResourceDimension RD;
2243 if (!HLSLResourceDimensionAttr::ConvertStrToResourceDimension(Identifier,
2245 Diag(ArgLoc, diag::warn_attribute_type_not_supported)
2246 <<
"ResourceDimension" << Identifier;
2249 A = HLSLResourceDimensionAttr::Create(
getASTContext(), RD, ACI);
2253 case ParsedAttr::AT_HLSLROV:
2257 case ParsedAttr::AT_HLSLRawBuffer:
2261 case ParsedAttr::AT_HLSLIsCounter:
2265 case ParsedAttr::AT_HLSLContainedType: {
2267 Diag(AL.
getLoc(), diag::err_attribute_wrong_number_arguments) << AL << 1;
2273 assert(TSI &&
"no type source info for attribute argument");
2275 diag::err_incomplete_type))
2277 A = HLSLContainedTypeAttr::Create(
getASTContext(), TSI, ACI);
2282 llvm_unreachable(
"unhandled HLSL attribute");
2285 HLSLResourcesTypeAttrs.emplace_back(A);
2291 if (!HLSLResourcesTypeAttrs.size())
2297 HLSLResourcesTypeAttrs, QT, &LocInfo)) {
2298 const HLSLAttributedResourceType *RT =
2305 LocsForHLSLAttributedResources.insert(std::pair(RT, LocInfo));
2307 HLSLResourcesTypeAttrs.clear();
2315 auto I = LocsForHLSLAttributedResources.find(RT);
2316 if (I != LocsForHLSLAttributedResources.end()) {
2317 LocInfo = I->second;
2318 LocsForHLSLAttributedResources.erase(I);
2327void SemaHLSL::collectResourceBindingsOnUserRecordDecl(
const VarDecl *VD,
2328 const RecordType *RT) {
2329 const RecordDecl *RD = RT->getDecl()->getDefinitionOrSelf();
2336 "incomplete arrays inside user defined types are not supported");
2345 if (
const HLSLAttributedResourceType *AttrResType =
2346 HLSLAttributedResourceType::findHandleTypeOnResource(Ty)) {
2351 Bindings.addDeclBindingInfo(VD, RC);
2352 }
else if (
const RecordType *RT = dyn_cast<RecordType>(Ty)) {
2358 collectResourceBindingsOnUserRecordDecl(VD, RT);
2370 bool SpecifiedSpace) {
2371 int RegTypeNum =
static_cast<int>(RegType);
2374 if (D->
hasAttr<HLSLGroupSharedAddressSpaceAttr>()) {
2375 S.
Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
2380 if (
HLSLBufferDecl *CBufferOrTBuffer = dyn_cast<HLSLBufferDecl>(D)) {
2381 ResourceClass RC = CBufferOrTBuffer->isCBuffer() ? ResourceClass::CBuffer
2382 : ResourceClass::SRV;
2392 assert(
isa<VarDecl>(D) &&
"D is expected to be VarDecl or HLSLBufferDecl");
2396 if (
const HLSLAttributedResourceType *AttrResType =
2397 HLSLAttributedResourceType::findHandleTypeOnResource(
2414 if (SpecifiedSpace && !DeclaredInCOrTBuffer)
2415 S.
Diag(ArgLoc, diag::err_hlsl_space_on_global_constant);
2420 if (RegType == RegisterType::CBuffer)
2421 S.
Diag(ArgLoc, diag::warn_hlsl_deprecated_register_type_b);
2422 else if (RegType != RegisterType::C)
2423 S.
Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
2427 if (RegType == RegisterType::C)
2428 S.
Diag(ArgLoc, diag::warn_hlsl_register_type_c_packoffset);
2430 S.
Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
2440 S.
Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
2448 bool RegisterTypesDetected[5] = {
false};
2449 RegisterTypesDetected[
static_cast<int>(regType)] =
true;
2452 if (HLSLResourceBindingAttr *
attr =
2453 dyn_cast<HLSLResourceBindingAttr>(*it)) {
2456 if (RegisterTypesDetected[
static_cast<int>(otherRegType)]) {
2457 int otherRegTypeNum =
static_cast<int>(otherRegType);
2459 diag::err_hlsl_duplicate_register_annotation)
2463 RegisterTypesDetected[
static_cast<int>(otherRegType)] =
true;
2471 bool SpecifiedSpace) {
2476 "expecting VarDecl or HLSLBufferDecl");
2488 const uint64_t &Limit,
2491 uint64_t ArrayCount = 1) {
2496 if (StartSlot > Limit)
2500 if (
const auto *AT = dyn_cast<ArrayType>(T)) {
2503 if (
const auto *CAT = dyn_cast<ConstantArrayType>(AT))
2504 Count = CAT->
getSize().getZExtValue();
2508 ArrayCount * Count);
2512 if (
auto ResTy = dyn_cast<HLSLAttributedResourceType>(T)) {
2515 if (ResTy->getAttrs().ResourceClass != ResClass)
2519 uint64_t EndSlot = StartSlot + ArrayCount - 1;
2520 if (EndSlot > Limit)
2524 StartSlot = EndSlot + 1;
2529 if (
const auto *RT = dyn_cast<RecordType>(T)) {
2532 if (
const auto *CXXRD = dyn_cast<CXXRecordDecl>(RD)) {
2535 ResClass, Ctx, ArrayCount))
2542 ResClass, Ctx, ArrayCount))
2556 const uint64_t Limit = UINT32_MAX;
2557 if (SlotNum > Limit)
2562 if (RegTy == RegisterType::C || RegTy == RegisterType::I)
2565 if (
VarDecl *VD = dyn_cast<VarDecl>(TheDecl)) {
2566 uint64_t BaseSlot = SlotNum;
2574 return (BaseSlot > Limit);
2581 return (SlotNum > Limit);
2584 llvm_unreachable(
"unexpected decl type");
2588 if (
VarDecl *VD = dyn_cast<VarDecl>(TheDecl)) {
2590 if (
const auto *IAT = dyn_cast<IncompleteArrayType>(Ty))
2591 Ty = IAT->getElementType();
2593 diag::err_incomplete_type))
2597 StringRef Slot =
"";
2598 StringRef Space =
"";
2602 Diag(AL.
getLoc(), diag::err_attribute_argument_type)
2612 Diag(AL.
getLoc(), diag::err_attribute_argument_type)
2618 SpaceLoc = Loc->
getLoc();
2621 if (Str.starts_with(
"space")) {
2623 SpaceLoc = Loc->
getLoc();
2632 std::optional<unsigned> SlotNum;
2633 unsigned SpaceNum = 0;
2636 if (!Slot.empty()) {
2638 Diag(SlotLoc, diag::err_hlsl_binding_type_invalid) << Slot.substr(0, 1);
2641 if (RegType == RegisterType::I) {
2642 Diag(SlotLoc, diag::warn_hlsl_deprecated_register_type_i);
2645 const StringRef SlotNumStr = Slot.substr(1);
2650 if (SlotNumStr.getAsInteger(10, N)) {
2651 Diag(SlotLoc, diag::err_hlsl_unsupported_register_number);
2659 Diag(SlotLoc, diag::err_hlsl_register_number_too_large);
2668 if (!Space.starts_with(
"space")) {
2669 Diag(SpaceLoc, diag::err_hlsl_expected_space) << Space;
2672 StringRef SpaceNumStr = Space.substr(5);
2673 if (SpaceNumStr.getAsInteger(10, SpaceNum)) {
2674 Diag(SpaceLoc, diag::err_hlsl_expected_space) << Space;
2679 if (SlotNum.has_value())
2684 HLSLResourceBindingAttr *NewAttr =
2685 HLSLResourceBindingAttr::Create(
getASTContext(), Slot, Space, AL);
2687 NewAttr->setBinding(RegType, SlotNum, SpaceNum);
2711 if (
auto *VD = dyn_cast<ValueDecl>(D))
2713 else if (
auto *TD = dyn_cast<TypedefNameDecl>(D))
2714 Ty = TD->getUnderlyingType();
2723 SemaRef.
Diag(Loc, diag::err_hlsl_matrix_layout_non_matrix) << AttrName;
2730 SemaRef.
Diag(Loc, diag::err_hlsl_matrix_layout_non_matrix) << AttrName;
2741 if (
const auto *Existing = D->
getAttr<HLSLMatrixLayoutAttr>()) {
2743 Diag(AL.
getLoc(), diag::err_hlsl_matrix_layout_conflict)
2745 Diag(Existing->getLoc(), diag::note_conflicting_attribute);
2747 Diag(AL.
getLoc(), diag::warn_duplicate_attribute_exact)
2749 Diag(Existing->getLoc(), diag::note_previous_attribute);
2758 Decl *D,
const HLSLMatrixLayoutAttr *
Attr) {
2805 llvm::DenseMap<const FunctionDecl *, unsigned> ScannedDecls;
2809 llvm::Triple::EnvironmentType CurrentShaderEnvironment;
2810 unsigned CurrentShaderStageBit;
2815 bool ReportOnlyShaderStageIssues;
2818 void SetShaderStageContext(llvm::Triple::EnvironmentType ShaderType) {
2819 static_assert(
sizeof(
unsigned) >= 4);
2820 assert(HLSLShaderAttr::isValidShaderType(ShaderType));
2821 assert((
unsigned)(ShaderType - llvm::Triple::Pixel) < 31 &&
2822 "ShaderType is too big for this bitmap");
2825 unsigned bitmapIndex = ShaderType - llvm::Triple::Pixel;
2826 CurrentShaderEnvironment = ShaderType;
2827 CurrentShaderStageBit = (1 << bitmapIndex);
2830 void SetUnknownShaderStageContext() {
2831 CurrentShaderEnvironment = llvm::Triple::UnknownEnvironment;
2832 CurrentShaderStageBit = (1 << 31);
2835 llvm::Triple::EnvironmentType GetCurrentShaderEnvironment()
const {
2836 return CurrentShaderEnvironment;
2839 bool InUnknownShaderStageContext()
const {
2840 return CurrentShaderEnvironment == llvm::Triple::UnknownEnvironment;
2844 void AddToScannedFunctions(
const FunctionDecl *FD) {
2845 unsigned &ScannedStages = ScannedDecls[FD];
2846 ScannedStages |= CurrentShaderStageBit;
2849 unsigned GetScannedStages(
const FunctionDecl *FD) {
return ScannedDecls[FD]; }
2851 bool WasAlreadyScannedInCurrentStage(
const FunctionDecl *FD) {
2852 return WasAlreadyScannedInCurrentStage(GetScannedStages(FD));
2855 bool WasAlreadyScannedInCurrentStage(
unsigned ScannerStages) {
2856 return ScannerStages & CurrentShaderStageBit;
2859 static bool NeverBeenScanned(
unsigned ScannedStages) {
2860 return ScannedStages == 0;
2864 void HandleFunctionOrMethodRef(FunctionDecl *FD, Expr *RefExpr);
2865 void CheckDeclAvailability(NamedDecl *D,
const AvailabilityAttr *AA,
2867 const AvailabilityAttr *FindAvailabilityAttr(
const Decl *D);
2868 bool HasMatchingEnvironmentOrNone(
const AvailabilityAttr *AA);
2871 DiagnoseHLSLAvailability(Sema &SemaRef)
2873 CurrentShaderEnvironment(llvm::Triple::UnknownEnvironment),
2874 CurrentShaderStageBit(0), ReportOnlyShaderStageIssues(
false) {}
2877 void RunOnTranslationUnit(
const TranslationUnitDecl *TU);
2878 void RunOnFunction(
const FunctionDecl *FD);
2880 bool VisitDeclRefExpr(DeclRefExpr *DRE)
override {
2881 FunctionDecl *FD = llvm::dyn_cast<FunctionDecl>(DRE->
getDecl());
2883 HandleFunctionOrMethodRef(FD, DRE);
2887 bool VisitMemberExpr(MemberExpr *ME)
override {
2888 FunctionDecl *FD = llvm::dyn_cast<FunctionDecl>(ME->
getMemberDecl());
2890 HandleFunctionOrMethodRef(FD, ME);
2895void DiagnoseHLSLAvailability::HandleFunctionOrMethodRef(
FunctionDecl *FD,
2898 "expected DeclRefExpr or MemberExpr");
2902 if (FD->
hasBody(FDWithBody)) {
2903 if (!WasAlreadyScannedInCurrentStage(FDWithBody))
2904 DeclsToScan.push_back(FDWithBody);
2909 const AvailabilityAttr *AA = FindAvailabilityAttr(FD);
2911 CheckDeclAvailability(
2915void DiagnoseHLSLAvailability::RunOnTranslationUnit(
2924 DeclContextsToScan.push_back(TU);
2926 while (!DeclContextsToScan.empty()) {
2927 const DeclContext *DC = DeclContextsToScan.pop_back_val();
2928 for (
auto &D : DC->
decls()) {
2935 if (llvm::dyn_cast<NamespaceDecl>(D) || llvm::dyn_cast<ExportDecl>(D)) {
2936 DeclContextsToScan.push_back(llvm::dyn_cast<DeclContext>(D));
2941 const FunctionDecl *FD = llvm::dyn_cast<FunctionDecl>(D);
2946 if (HLSLShaderAttr *ShaderAttr = FD->
getAttr<HLSLShaderAttr>()) {
2947 SetShaderStageContext(ShaderAttr->getType());
2956 for (
const auto *Redecl : FD->
redecls()) {
2957 if (Redecl->isInExportDeclContext()) {
2964 SetUnknownShaderStageContext();
2972void DiagnoseHLSLAvailability::RunOnFunction(
const FunctionDecl *FD) {
2973 assert(DeclsToScan.empty() &&
"DeclsToScan should be empty");
2974 DeclsToScan.push_back(FD);
2976 while (!DeclsToScan.empty()) {
2984 const unsigned ScannedStages = GetScannedStages(FD);
2985 if (WasAlreadyScannedInCurrentStage(ScannedStages))
2988 ReportOnlyShaderStageIssues = !NeverBeenScanned(ScannedStages);
2990 AddToScannedFunctions(FD);
2995bool DiagnoseHLSLAvailability::HasMatchingEnvironmentOrNone(
2996 const AvailabilityAttr *AA) {
3001 llvm::Triple::EnvironmentType CurrentEnv = GetCurrentShaderEnvironment();
3002 if (CurrentEnv == llvm::Triple::UnknownEnvironment)
3005 llvm::Triple::EnvironmentType AttrEnv =
3006 AvailabilityAttr::getEnvironmentType(IIEnvironment->
getName());
3008 return CurrentEnv == AttrEnv;
3011const AvailabilityAttr *
3012DiagnoseHLSLAvailability::FindAvailabilityAttr(
const Decl *D) {
3013 AvailabilityAttr
const *PartialMatch =
nullptr;
3017 for (
const auto *A : D->
attrs()) {
3018 if (
const auto *Avail = dyn_cast<AvailabilityAttr>(A)) {
3019 const AvailabilityAttr *EffectiveAvail = Avail->getEffectiveAttr();
3020 StringRef AttrPlatform = EffectiveAvail->getPlatform()->getName();
3021 StringRef TargetPlatform =
3025 if (AttrPlatform == TargetPlatform) {
3027 if (HasMatchingEnvironmentOrNone(EffectiveAvail))
3029 PartialMatch = Avail;
3033 return PartialMatch;
3038void DiagnoseHLSLAvailability::CheckDeclAvailability(
NamedDecl *D,
3039 const AvailabilityAttr *AA,
3058 if (ReportOnlyShaderStageIssues)
3064 if (InUnknownShaderStageContext())
3069 bool EnvironmentMatches = HasMatchingEnvironmentOrNone(AA);
3070 VersionTuple Introduced = AA->getIntroduced();
3079 llvm::StringRef PlatformName(
3082 llvm::StringRef CurrentEnvStr =
3083 llvm::Triple::getEnvironmentTypeName(GetCurrentShaderEnvironment());
3085 llvm::StringRef AttrEnvStr =
3086 AA->getEnvironment() ? AA->getEnvironment()->getName() :
"";
3087 bool UseEnvironment = !AttrEnvStr.empty();
3089 if (EnvironmentMatches) {
3090 SemaRef.
Diag(
Range.getBegin(), diag::warn_hlsl_availability)
3091 <<
Range << D << PlatformName << Introduced.getAsString()
3092 << UseEnvironment << CurrentEnvStr;
3094 SemaRef.
Diag(
Range.getBegin(), diag::warn_hlsl_availability_unavailable)
3098 SemaRef.
Diag(D->
getLocation(), diag::note_partial_availability_specified_here)
3099 << D << PlatformName << Introduced.getAsString()
3101 << UseEnvironment << AttrEnvStr << CurrentEnvStr;
3108 if (!DefaultCBufferDecls.empty()) {
3111 DefaultCBufferDecls);
3114 SemaRef.getCurLexicalContext()->addDecl(DefaultCBuffer);
3118 for (
const Decl *VD : DefaultCBufferDecls) {
3119 const HLSLResourceBindingAttr *RBA =
3120 VD->
getAttr<HLSLResourceBindingAttr>();
3121 if (RBA && RBA->hasRegisterSlot() &&
3122 RBA->getRegisterType() == HLSLResourceBindingAttr::RegisterType::C) {
3129 SemaRef.Consumer.HandleTopLevelDecl(DG);
3131 diagnoseAvailabilityViolations(TU);
3140 "expected member expr to have resource record type or array of them");
3146 const Expr *NonConstIndexExpr =
nullptr;
3149 if (
const DeclRefExpr *DRE = dyn_cast<DeclRefExpr>(E)) {
3150 if (!NonConstIndexExpr)
3158 diag::err_hlsl_resource_member_array_access_not_constant);
3162 if (
const auto *ASE = dyn_cast<ArraySubscriptExpr>(E)) {
3163 const Expr *IdxExpr = ASE->getIdx();
3165 NonConstIndexExpr = IdxExpr;
3167 }
else if (
const auto *SubME = dyn_cast<MemberExpr>(E)) {
3168 E = SubME->getBase();
3169 }
else if (
const auto *ICE = dyn_cast<ImplicitCastExpr>(E)) {
3170 E = ICE->getSubExpr();
3172 llvm_unreachable(
"unexpected expr type in resource member access");
3181 SemaRef.Context.getCanonicalType(
SemaRef.Context.getAddrSpaceQualType(
3184 SemaRef.Context.getLValueReferenceType(AddrSpaceType));
3187 SemaRef.Context.DeclarationNames.getCXXConversionFunctionName(
3191 [[maybe_unused]]
bool LookupSucceeded =
3192 SemaRef.LookupQualifiedName(ConvR, RD);
3193 assert(LookupSucceeded);
3202std::optional<ExprResult>
3205 const HLSLAttributedResourceType *ResTy =
3206 HLSLAttributedResourceType::findHandleTypeOnResource(
3207 BaseType.getTypePtr());
3209 ResTy->getAttrs().ResourceClass != llvm::dxil::ResourceClass::CBuffer)
3210 return std::nullopt;
3212 QualType TemplateType = ResTy->getContainedType();
3216 assert(NamedConversionDecl &&
3217 "Could not find conversion function for ConstantBuffer.");
3218 auto *ConversionDecl =
3221 return SemaRef.BuildCXXMemberCallExpr(BaseExpr.
get(), NamedConversionDecl,
3233 TI.
getTriple().getEnvironment() != llvm::Triple::EnvironmentType::Library)
3236 DiagnoseHLSLAvailability(
SemaRef).RunOnTranslationUnit(TU);
3243 for (
unsigned I = 1, N = TheCall->
getNumArgs(); I < N; ++I) {
3246 S->
Diag(TheCall->
getBeginLoc(), diag::err_vec_builtin_incompatible_vector)
3271 for (
unsigned I = 0; I < TheCall->
getNumArgs(); ++I) {
3286 if (!BaseType->isFloat32Type())
3287 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3288 << ArgOrdinal << 5 << 0
3300 if (!BaseType->isHalfType() && !BaseType->isFloat32Type())
3301 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3302 << ArgOrdinal << 5 << 0
3316 if (!BaseType->isDoubleType()) {
3319 return S->
Diag(Loc, diag::err_builtin_requires_double_type)
3320 << ArgOrdinal << PassedType;
3327 unsigned ArgIndex) {
3328 auto *Arg = TheCall->
getArg(ArgIndex);
3330 if (Arg->IgnoreCasts()->isModifiableLvalue(S->
Context, &OrigLoc) ==
3333 S->
Diag(OrigLoc, diag::error_hlsl_inout_lvalue) << Arg << 0;
3343 if (VecTy->getElementType()->isDoubleType())
3344 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3345 << ArgOrdinal << 1 << 0 << 1
3355 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3356 << ArgOrdinal << 5 << 1
3365 if (VecTy->getElementType()->isUnsignedIntegerType())
3368 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3369 << ArgOrdinal << 4 << 3 << 0
3378 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3379 << ArgOrdinal << 5 << 3
3385 unsigned ArgOrdinal,
unsigned Width) {
3388 ArgTy = VTy->getElementType();
3390 uint64_t ElementBitCount =
3392 if (ElementBitCount != Width) {
3394 diag::err_integer_incorrect_bit_count)
3395 << Width << ElementBitCount;
3412 unsigned ArgIndex) {
3421 diag::err_typecheck_expect_scalar_or_vector)
3422 << ArgType << Scalar;
3429 QualType Scalar,
unsigned ArgIndex) {
3440 if (
const auto *VTy = ArgType->getAs<
VectorType>()) {
3453 diag::err_typecheck_expect_scalar_or_vector_or_matrix)
3454 << ArgType << Scalar;
3459 unsigned ArgIndex) {
3464 if (!(ArgType->isScalarType() ||
3465 (VTy && VTy->getElementType()->isScalarType()))) {
3467 diag::err_typecheck_expect_any_scalar_or_vector)
3477 unsigned ArgIndex) {
3479 assert(ArgIndex < TheCall->getNumArgs());
3487 diag::err_typecheck_expect_any_scalar_or_vector)
3512 diag::err_typecheck_call_different_arg_types)
3531 Arg1ScalarTy = VTy->getElementType();
3535 Arg2ScalarTy = VTy->getElementType();
3538 S->
Diag(Arg1->
getBeginLoc(), diag::err_hlsl_builtin_scalar_vector_mismatch)
3539 << 1 << TheCall->
getCallee() << Arg1Ty << Arg2Ty;
3549 if (Arg1Length > 0 && Arg0Length != Arg1Length) {
3551 diag::err_typecheck_vector_lengths_not_equal)
3557 if (Arg2Length > 0 && Arg0Length != Arg2Length) {
3559 diag::err_typecheck_vector_lengths_not_equal)
3571 assert(TheCall->
getNumArgs() > IndexArgIndex &&
"Index argument missing");
3574 unsigned int ActualDim = 1;
3576 ActualDim = VTy->getNumElements();
3577 IndexTy = VTy->getElementType();
3581 diag::err_typecheck_expect_int)
3587 const HLSLAttributedResourceType *ResTy =
3589 assert(ResTy &&
"Resource argument must be a resource");
3590 HLSLAttributedResourceType::Attributes ResAttrs = ResTy->getAttrs();
3592 unsigned int ExpectedDim = 1;
3593 if (ResAttrs.ResourceDimension != llvm::dxil::ResourceDimension::Unknown)
3596 if (ActualDim != ExpectedDim) {
3598 diag::err_hlsl_builtin_resource_coordinate_dimension_mismatch)
3609 llvm::function_ref<
bool(
const HLSLAttributedResourceType *ResType)> Check =
3613 const HLSLAttributedResourceType *ResTy =
3617 diag::err_typecheck_expect_hlsl_resource)
3621 if (Check && Check(ResTy)) {
3623 diag::err_invalid_hlsl_resource_type)
3631 QualType BaseType,
unsigned ExpectedCount,
3633 unsigned PassedCount = 1;
3635 PassedCount = VecTy->getNumElements();
3637 if (PassedCount != ExpectedCount) {
3640 S->
Diag(Loc, diag::err_typecheck_convert_incompatible)
3652 [](
const HLSLAttributedResourceType *ResType) {
3653 return ResType->getAttrs().ResourceDimension ==
3654 llvm::dxil::ResourceDimension::Unknown;
3660 [](
const HLSLAttributedResourceType *ResType) {
3661 return ResType->getAttrs().ResourceClass !=
3662 llvm::hlsl::ResourceClass::Sampler;
3670 unsigned ExpectedDim =
3698 unsigned NextIdx = 3;
3704 diag::err_typecheck_convert_incompatible)
3712 Expr *ComponentArg = TheCall->
getArg(NextIdx);
3716 diag::err_typecheck_convert_incompatible)
3723 std::optional<llvm::APSInt> ComponentOpt =
3726 int64_t ComponentVal = ComponentOpt->getSExtValue();
3727 if (ComponentVal != 0) {
3730 assert(ComponentVal >= 0 && ComponentVal <= 3 &&
3731 "The component is not in the expected range.");
3733 diag::err_hlsl_gathercmp_invalid_component)
3743 const HLSLAttributedResourceType *ResourceTy =
3746 unsigned ExpectedDim =
3755 assert(ResourceTy->hasContainedType() &&
3756 "Expecting a contained type for resource with a dimension "
3758 QualType ReturnType = ResourceTy->getContainedType();
3762 S.
Diag(TheCall->
getBeginLoc(), diag::err_hlsl_samplecmp_requires_float);
3768 ReturnType = VecTy->getElementType();
3781 [](
const HLSLAttributedResourceType *ResType) {
3782 return ResType->getAttrs().ResourceDimension ==
3783 llvm::dxil::ResourceDimension::Unknown;
3791 unsigned ExpectedDim =
3800 EltTy = VTy->getElementType();
3815 TheCall->
setType(ResourceTy->getContainedType());
3820 unsigned MinArgs, MaxArgs;
3848 const HLSLAttributedResourceType *ResourceTy =
3850 unsigned ExpectedDim =
3853 unsigned NextIdx = 3;
3862 diag::err_typecheck_convert_incompatible)
3897 diag::err_typecheck_convert_incompatible)
3903 assert(ResourceTy->hasContainedType() &&
3904 "Expecting a contained type for resource with a dimension "
3906 QualType ReturnType = ResourceTy->getContainedType();
3909 S.
Diag(TheCall->
getBeginLoc(), diag::err_hlsl_samplecmp_requires_float);
3922 switch (BuiltinID) {
3923 case Builtin::BI__builtin_hlsl_adduint64: {
3924 if (
SemaRef.checkArgCount(TheCall, 2))
3938 if (NumElementsArg != 2 && NumElementsArg != 4) {
3940 << 1 << 64 << NumElementsArg * 32;
3954 case Builtin::BI__builtin_hlsl_resource_getpointer: {
3955 if (
SemaRef.checkArgCountRange(TheCall, 1, 2) ||
3962 QualType ContainedTy = ResourceTy->getContainedType();
3963 auto ReturnType =
SemaRef.Context.getAddrSpaceQualType(
3966 ReturnType =
SemaRef.Context.getPointerType(ReturnType);
3971 case Builtin::BI__builtin_hlsl_resource_getpointer_typed: {
3972 if (
SemaRef.checkArgCount(TheCall, 3) ||
3979 "expected pointer type for second argument");
3986 diag::err_invalid_use_of_array_type);
3990 auto ReturnType =
SemaRef.Context.getAddrSpaceQualType(
3993 ReturnType =
SemaRef.Context.getPointerType(ReturnType);
3998 case Builtin::BI__builtin_hlsl_resource_load_with_status: {
3999 if (
SemaRef.checkArgCount(TheCall, 3) ||
4002 SemaRef.getASTContext().UnsignedIntTy) ||
4004 SemaRef.getASTContext().UnsignedIntTy) ||
4010 QualType ReturnType = ResourceTy->getContainedType();
4015 case Builtin::BI__builtin_hlsl_resource_load_with_status_typed: {
4016 if (
SemaRef.checkArgCount(TheCall, 4) ||
4019 SemaRef.getASTContext().UnsignedIntTy) ||
4021 SemaRef.getASTContext().UnsignedIntTy) ||
4027 "expected pointer type for second argument");
4034 diag::err_invalid_use_of_array_type);
4040 case Builtin::BI__builtin_hlsl_resource_load_level:
4042 case Builtin::BI__builtin_hlsl_resource_sample:
4044 case Builtin::BI__builtin_hlsl_resource_sample_bias:
4046 case Builtin::BI__builtin_hlsl_resource_sample_grad:
4048 case Builtin::BI__builtin_hlsl_resource_sample_level:
4050 case Builtin::BI__builtin_hlsl_resource_sample_cmp:
4052 case Builtin::BI__builtin_hlsl_resource_sample_cmp_level_zero:
4054 case Builtin::BI__builtin_hlsl_resource_calculate_lod:
4055 case Builtin::BI__builtin_hlsl_resource_calculate_lod_unclamped:
4057 case Builtin::BI__builtin_hlsl_resource_gather:
4059 case Builtin::BI__builtin_hlsl_resource_gather_cmp:
4061 case Builtin::BI__builtin_hlsl_resource_uninitializedhandle: {
4062 assert(TheCall->
getNumArgs() == 1 &&
"expected 1 arg");
4068 case Builtin::BI__builtin_hlsl_resource_handlefrombinding: {
4069 assert(TheCall->
getNumArgs() == 6 &&
"expected 6 args");
4075 case Builtin::BI__builtin_hlsl_resource_handlefromimplicitbinding: {
4076 assert(TheCall->
getNumArgs() == 6 &&
"expected 6 args");
4082 case Builtin::BI__builtin_hlsl_resource_counterhandlefromimplicitbinding: {
4083 assert(TheCall->
getNumArgs() == 3 &&
"expected 3 args");
4086 auto *MainResType = MainHandleTy->
getAs<HLSLAttributedResourceType>();
4087 auto MainAttrs = MainResType->getAttrs();
4088 assert(!MainAttrs.IsCounter &&
"cannot create a counter from a counter");
4089 MainAttrs.IsCounter =
true;
4091 MainResType->getWrappedType(), MainResType->getContainedType(),
4095 TheCall->
setType(CounterHandleTy);
4098 case Builtin::BI__builtin_hlsl_and:
4099 case Builtin::BI__builtin_hlsl_or: {
4100 if (
SemaRef.checkArgCount(TheCall, 2))
4114 case Builtin::BI__builtin_hlsl_all:
4115 case Builtin::BI__builtin_hlsl_any: {
4116 if (
SemaRef.checkArgCount(TheCall, 1))
4122 case Builtin::BI__builtin_hlsl_asdouble: {
4123 if (
SemaRef.checkArgCount(TheCall, 2))
4127 SemaRef.Context.UnsignedIntTy,
4132 SemaRef.Context.UnsignedIntTy,
4141 case Builtin::BI__builtin_hlsl_elementwise_clamp: {
4142 if (
SemaRef.BuiltinElementwiseTernaryMath(
4148 case Builtin::BI__builtin_hlsl_dot: {
4150 if (
SemaRef.BuiltinVectorToScalarMath(TheCall))
4156 case Builtin::BI__builtin_hlsl_elementwise_firstbithigh:
4157 case Builtin::BI__builtin_hlsl_elementwise_firstbitlow: {
4158 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
4168 EltTy = VecTy->getElementType();
4169 ResTy =
SemaRef.Context.getExtVectorType(ResTy, VecTy->getNumElements());
4182 case Builtin::BI__builtin_hlsl_select: {
4183 if (
SemaRef.checkArgCount(TheCall, 3))
4191 if (VTy && VTy->getElementType()->isBooleanType() &&
4196 case Builtin::BI__builtin_hlsl_elementwise_saturate:
4197 case Builtin::BI__builtin_hlsl_elementwise_rcp: {
4198 if (
SemaRef.checkArgCount(TheCall, 1))
4204 diag::err_builtin_invalid_arg_type)
4207 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
4211 case Builtin::BI__builtin_hlsl_elementwise_degrees:
4212 case Builtin::BI__builtin_hlsl_elementwise_radians:
4213 case Builtin::BI__builtin_hlsl_elementwise_rsqrt:
4214 case Builtin::BI__builtin_hlsl_elementwise_frac:
4215 case Builtin::BI__builtin_hlsl_elementwise_ddx_coarse:
4216 case Builtin::BI__builtin_hlsl_elementwise_ddy_coarse:
4217 case Builtin::BI__builtin_hlsl_elementwise_ddx_fine:
4218 case Builtin::BI__builtin_hlsl_elementwise_ddy_fine: {
4219 if (
SemaRef.checkArgCount(TheCall, 1))
4224 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
4228 case Builtin::BI__builtin_hlsl_elementwise_isinf:
4229 case Builtin::BI__builtin_hlsl_elementwise_isnan: {
4230 if (
SemaRef.checkArgCount(TheCall, 1))
4235 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
4240 case Builtin::BI__builtin_hlsl_lerp: {
4241 if (
SemaRef.checkArgCount(TheCall, 3))
4248 if (
SemaRef.BuiltinElementwiseTernaryMath(TheCall))
4252 case Builtin::BI__builtin_hlsl_mad: {
4253 if (
SemaRef.BuiltinElementwiseTernaryMath(
4259 case Builtin::BI__builtin_hlsl_mul: {
4260 if (
SemaRef.checkArgCount(TheCall, 2))
4269 if (
const auto *VTy = T->getAs<
VectorType>())
4270 return VTy->getElementType();
4272 return MTy->getElementType();
4276 QualType EltTy0 = getElemType(Ty0);
4285 if (IsVec0 && IsMat1) {
4288 }
else if (IsMat0 && IsVec1) {
4292 assert(IsMat0 && IsMat1);
4302 case Builtin::BI__builtin_hlsl_normalize: {
4303 if (
SemaRef.checkArgCount(TheCall, 1))
4314 case Builtin::BI__builtin_elementwise_fma: {
4315 if (
SemaRef.checkArgCount(TheCall, 3) ||
4330 case Builtin::BI__builtin_hlsl_transpose: {
4331 if (
SemaRef.checkArgCount(TheCall, 1))
4340 << 1 << 3 << 0 << 0 << ArgTy;
4345 MatTy->getElementType(), MatTy->getNumColumns(), MatTy->getNumRows());
4349 case Builtin::BI__builtin_hlsl_elementwise_sign: {
4350 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
4358 case Builtin::BI__builtin_hlsl_step: {
4359 if (
SemaRef.checkArgCount(TheCall, 2))
4371 case Builtin::BI__builtin_hlsl_wave_active_all_equal: {
4372 if (
SemaRef.checkArgCount(TheCall, 1))
4386 unsigned NumElts = VecTy->getNumElements();
4396 case Builtin::BI__builtin_hlsl_wave_active_max:
4397 case Builtin::BI__builtin_hlsl_wave_active_min:
4398 case Builtin::BI__builtin_hlsl_wave_active_sum:
4399 case Builtin::BI__builtin_hlsl_wave_active_product: {
4400 if (
SemaRef.checkArgCount(TheCall, 1))
4413 case Builtin::BI__builtin_hlsl_wave_active_bit_or:
4414 case Builtin::BI__builtin_hlsl_wave_active_bit_xor:
4415 case Builtin::BI__builtin_hlsl_wave_active_bit_and: {
4416 if (
SemaRef.checkArgCount(TheCall, 1))
4431 (VTy && VTy->getElementType()->isIntegerType()))) {
4433 diag::err_builtin_invalid_arg_type)
4434 << ArgTyExpr <<
SemaRef.Context.UnsignedIntTy << 1 << 0 << 0;
4444 case Builtin::BI__builtin_elementwise_bitreverse: {
4452 case Builtin::BI__builtin_hlsl_wave_prefix_count_bits: {
4453 if (
SemaRef.checkArgCount(TheCall, 1))
4458 if (!(
ArgType->isScalarType())) {
4460 diag::err_typecheck_expect_any_scalar_or_vector)
4465 if (!(
ArgType->isBooleanType())) {
4467 diag::err_typecheck_expect_any_scalar_or_vector)
4474 case Builtin::BI__builtin_hlsl_wave_read_lane_at: {
4475 if (
SemaRef.checkArgCount(TheCall, 2))
4483 diag::err_typecheck_convert_incompatible)
4484 << ArgTyIndex <<
SemaRef.Context.UnsignedIntTy << 1 << 0 << 0;
4497 case Builtin::BI__builtin_hlsl_wave_get_lane_index: {
4498 if (
SemaRef.checkArgCount(TheCall, 0))
4502 case Builtin::BI__builtin_hlsl_wave_prefix_sum:
4503 case Builtin::BI__builtin_hlsl_wave_prefix_product: {
4504 if (
SemaRef.checkArgCount(TheCall, 1))
4517 case Builtin::BI__builtin_hlsl_quad_read_across_x:
4518 case Builtin::BI__builtin_hlsl_quad_read_across_y: {
4519 if (
SemaRef.checkArgCount(TheCall, 1))
4531 case Builtin::BI__builtin_hlsl_elementwise_splitdouble: {
4532 if (
SemaRef.checkArgCount(TheCall, 3))
4547 case Builtin::BI__builtin_hlsl_elementwise_clip: {
4548 if (
SemaRef.checkArgCount(TheCall, 1))
4555 case Builtin::BI__builtin_elementwise_acos:
4556 case Builtin::BI__builtin_elementwise_asin:
4557 case Builtin::BI__builtin_elementwise_atan:
4558 case Builtin::BI__builtin_elementwise_atan2:
4559 case Builtin::BI__builtin_elementwise_ceil:
4560 case Builtin::BI__builtin_elementwise_cos:
4561 case Builtin::BI__builtin_elementwise_cosh:
4562 case Builtin::BI__builtin_elementwise_exp:
4563 case Builtin::BI__builtin_elementwise_exp2:
4564 case Builtin::BI__builtin_elementwise_exp10:
4565 case Builtin::BI__builtin_elementwise_floor:
4566 case Builtin::BI__builtin_elementwise_fmod:
4567 case Builtin::BI__builtin_elementwise_log:
4568 case Builtin::BI__builtin_elementwise_log2:
4569 case Builtin::BI__builtin_elementwise_log10:
4570 case Builtin::BI__builtin_elementwise_pow:
4571 case Builtin::BI__builtin_elementwise_roundeven:
4572 case Builtin::BI__builtin_elementwise_sin:
4573 case Builtin::BI__builtin_elementwise_sinh:
4574 case Builtin::BI__builtin_elementwise_sqrt:
4575 case Builtin::BI__builtin_elementwise_tan:
4576 case Builtin::BI__builtin_elementwise_tanh:
4577 case Builtin::BI__builtin_elementwise_trunc: {
4583 case Builtin::BI__builtin_hlsl_buffer_update_counter: {
4584 assert(TheCall->
getNumArgs() == 2 &&
"expected 2 args");
4585 auto checkResTy = [](
const HLSLAttributedResourceType *ResTy) ->
bool {
4586 return !(ResTy->getAttrs().ResourceClass == ResourceClass::UAV &&
4587 ResTy->getAttrs().RawBuffer && ResTy->hasContainedType());
4592 std::optional<llvm::APSInt> Offset =
4594 if (!Offset.has_value() ||
std::abs(Offset->getExtValue()) != 1) {
4596 diag::err_hlsl_expect_arg_const_int_one_or_neg_one)
4602 case Builtin::BI__builtin_hlsl_elementwise_f16tof32: {
4603 if (
SemaRef.checkArgCount(TheCall, 1))
4614 ArgTy = VTy->getElementType();
4617 diag::err_builtin_invalid_arg_type)
4626 case Builtin::BI__builtin_hlsl_elementwise_f32tof16: {
4627 if (
SemaRef.checkArgCount(TheCall, 1))
4642 WorkList.push_back(BaseTy);
4643 while (!WorkList.empty()) {
4644 QualType T = WorkList.pop_back_val();
4645 T = T.getCanonicalType().getUnqualifiedType();
4646 if (
const auto *AT = dyn_cast<ConstantArrayType>(T)) {
4654 for (uint64_t Ct = 0; Ct < AT->
getZExtSize(); ++Ct)
4655 llvm::append_range(List, ElementFields);
4660 if (
const auto *VT = dyn_cast<VectorType>(T)) {
4661 List.insert(List.end(), VT->getNumElements(), VT->getElementType());
4664 if (
const auto *MT = dyn_cast<ConstantMatrixType>(T)) {
4665 List.insert(List.end(), MT->getNumElementsFlattened(),
4666 MT->getElementType());
4669 if (
const auto *RD = T->getAsCXXRecordDecl()) {
4670 if (RD->isStandardLayout())
4671 RD = RD->getStandardLayoutBaseWithFields();
4675 if (RD->
isUnion() || !RD->isAggregate()) {
4681 for (
const auto *FD : RD->
fields())
4682 if (!FD->isUnnamedBitField())
4683 FieldTypes.push_back(FD->
getType());
4685 std::reverse(FieldTypes.begin(), FieldTypes.end());
4686 llvm::append_range(WorkList, FieldTypes);
4690 if (!RD->isStandardLayout()) {
4692 for (
const auto &
Base : RD->bases())
4693 FieldTypes.push_back(
Base.getType());
4694 std::reverse(FieldTypes.begin(), FieldTypes.end());
4695 llvm::append_range(WorkList, FieldTypes);
4717 if (
SemaRef.Context.getTypeSize(QT) / 8 > 16)
4723 int ArraySize = VT->getNumElements();
4728 QualType ElTy = VT->getElementType();
4732 if (
SemaRef.Context.getTypeSize(QT) / 8 > 16)
4748 if (
SemaRef.getASTContext().hasSameType(T1, T2))
4757 return llvm::equal(T1Types, T2Types,
4759 return SemaRef.IsLayoutCompatible(LHS, RHS);
4768 bool HadError =
false;
4770 for (
unsigned i = 0, e =
New->getNumParams(); i != e; ++i) {
4778 const auto *NDAttr = NewParam->
getAttr<HLSLParamModifierAttr>();
4779 unsigned NSpellingIdx = (NDAttr ? NDAttr->getSpellingListIndex() : 0);
4780 const auto *ODAttr = OldParam->
getAttr<HLSLParamModifierAttr>();
4781 unsigned OSpellingIdx = (ODAttr ? ODAttr->getSpellingListIndex() : 0);
4783 if (NSpellingIdx != OSpellingIdx) {
4785 diag::err_hlsl_param_qualifier_mismatch)
4786 << NDAttr << NewParam;
4802 if (
SemaRef.getASTContext().hasSameUnqualifiedType(SrcTy, DestTy))
4817 llvm_unreachable(
"HLSL doesn't support pointers.");
4820 llvm_unreachable(
"HLSL doesn't support complex types.");
4822 llvm_unreachable(
"HLSL doesn't support fixed point types.");
4824 llvm_unreachable(
"Should have returned before this");
4834 llvm_unreachable(
"HLSL doesn't support complex types.");
4836 llvm_unreachable(
"HLSL doesn't support fixed point types.");
4841 llvm_unreachable(
"HLSL doesn't support pointers.");
4843 llvm_unreachable(
"Should have returned before this");
4849 llvm_unreachable(
"HLSL doesn't support pointers.");
4852 llvm_unreachable(
"HLSL doesn't support fixed point types.");
4856 llvm_unreachable(
"HLSL doesn't support complex types.");
4859 llvm_unreachable(
"Unhandled scalar cast");
4880 !(SrcMatTy && SrcMatTy->getNumElementsFlattened() == 1))
4886 SrcTy = SrcMatTy->getElementType();
4891 for (
unsigned I = 0, Size = DestTypes.size(); I < Size; ++I) {
4892 if (DestTypes[I]->isUnionType())
4924 if (SrcTypes.size() < DestTypes.size())
4927 unsigned SrcSize = SrcTypes.size();
4928 unsigned DstSize = DestTypes.size();
4930 for (I = 0; I < DstSize && I < SrcSize; I++) {
4931 if (SrcTypes[I]->isUnionType() || DestTypes[I]->isUnionType())
4939 for (; I < SrcSize; I++) {
4940 if (SrcTypes[I]->isUnionType())
4947 assert(Param->hasAttr<HLSLParamModifierAttr>() &&
4948 "We should not get here without a parameter modifier expression");
4949 const auto *
Attr = Param->getAttr<HLSLParamModifierAttr>();
4956 << Arg << (IsInOut ? 1 : 0);
4962 QualType Ty = Param->getType().getNonLValueExprType(Ctx);
4969 << Arg << (IsInOut ? 1 : 0);
4981 SemaRef.PerformCopyInitialization(Entity, Param->getBeginLoc(), ArgOpV);
4987 auto *OpV =
new (Ctx)
4992 Res =
SemaRef.ActOnBinOp(
SemaRef.getCurScope(), Param->getBeginLoc(),
4993 tok::equal, ArgOpV, OpV);
5009 "Pointer and reference types cannot be inout or out parameters");
5010 Ty =
SemaRef.getASTContext().getLValueReferenceType(Ty);
5026 for (
const auto *FD : RD->
fields()) {
5030 assert(RD->getNumBases() <= 1 &&
5031 "HLSL doesn't support multiple inheritance");
5032 return RD->getNumBases()
5037 if (
const auto *AT = dyn_cast<ArrayType>(Ty)) {
5038 if (
const auto *CAT = dyn_cast<ConstantArrayType>(AT))
5050 bool IsVKPushConstant = IsVulkan && VD->
hasAttr<HLSLVkPushConstantAttr>();
5055 !VD->
hasAttr<HLSLVkConstantIdAttr>() && !IsVKPushConstant &&
5061 if (
Decl->getType().hasAddressSpace())
5064 if (
Decl->getType()->isDependentType())
5076 if (
Decl->
hasAttr<HLSLVkExtBuiltinOutputAttr>()) {
5090 llvm::Triple::Vulkan;
5091 if (IsVulkan &&
Decl->
hasAttr<HLSLVkPushConstantAttr>()) {
5092 if (HasDeclaredAPushConstant)
5098 HasDeclaredAPushConstant =
true;
5125class StructBindingContext {
5128 HLSLResourceBindingAttr *RegBindingsAttrs[4];
5129 unsigned RegBindingOffset[4];
5132 static_assert(
static_cast<unsigned>(RegisterType::SRV) == 0 &&
5133 static_cast<unsigned>(RegisterType::UAV) == 1 &&
5134 static_cast<unsigned>(RegisterType::CBuffer) == 2 &&
5135 static_cast<unsigned>(RegisterType::Sampler) == 3,
5136 "unexpected register type values");
5139 HLSLVkBindingAttr *VkBindingAttr;
5140 unsigned VkBindingOffset;
5145 StructBindingContext(
VarDecl *VD) {
5146 for (
unsigned i = 0; i < 4; ++i) {
5147 RegBindingsAttrs[i] =
nullptr;
5148 RegBindingOffset[i] = 0;
5150 VkBindingAttr =
nullptr;
5151 VkBindingOffset = 0;
5157 if (
auto *RBA = dyn_cast<HLSLResourceBindingAttr>(A)) {
5159 unsigned RegTypeIdx =
static_cast<unsigned>(RegType);
5162 RegBindingsAttrs[RegTypeIdx] = RBA;
5167 if (
auto *VBA = dyn_cast<HLSLVkBindingAttr>(A))
5168 VkBindingAttr = VBA;
5175 Attr *createBindingAttr(SemaHLSL &S, ASTContext &AST,
RegisterType RegType,
5176 unsigned Range,
bool HasCounter) {
5177 assert(
static_cast<unsigned>(RegType) < 4 &&
"unexpected register type");
5179 if (VkBindingAttr) {
5180 unsigned Offset = VkBindingOffset;
5181 VkBindingOffset +=
Range;
5182 return HLSLVkBindingAttr::CreateImplicit(
5183 AST, VkBindingAttr->getBinding() + Offset, VkBindingAttr->getSet(),
5184 VkBindingAttr->getRange());
5187 HLSLResourceBindingAttr *RBA =
5188 RegBindingsAttrs[
static_cast<unsigned>(RegType)];
5189 HLSLResourceBindingAttr *NewAttr =
nullptr;
5191 if (RBA && RBA->hasRegisterSlot()) {
5194 unsigned Offset = RegBindingOffset[
static_cast<unsigned>(RegType)];
5195 RegBindingOffset[
static_cast<unsigned>(RegType)] += Range;
5197 unsigned NewSlotNumber = RBA->getSlotNumber() + Offset;
5198 StringRef NewSlotNumberStr =
5200 NewAttr = HLSLResourceBindingAttr::CreateImplicit(
5201 AST, NewSlotNumberStr, RBA->getSpace(), RBA->getRange());
5202 NewAttr->setBinding(RegType, NewSlotNumber, RBA->getSpaceNumber());
5206 NewAttr = HLSLResourceBindingAttr::CreateImplicit(AST,
"",
"0", {});
5207 NewAttr->setBinding(RegType, std::nullopt,
5208 RBA ? RBA->getSpaceNumber() : 0);
5212 NewAttr->setImplicitCounterBindingOrderID(
5221static void createGlobalResourceDeclForStruct(
5223 QualType ResTy, StructBindingContext &BindingCtx) {
5225 "expected resource type or array of resources");
5236 while (
const auto *AT = dyn_cast<ArrayType>(SingleResTy)) {
5237 const auto *CAT = dyn_cast<ConstantArrayType>(AT);
5242 const HLSLAttributedResourceType *ResHandleTy =
5243 HLSLAttributedResourceType::findHandleTypeOnResource(SingleResTy);
5247 Attr *BindingAttr = BindingCtx.createBindingAttr(
5249 ResDecl->
addAttr(BindingAttr);
5250 ResDecl->
addAttr(InternalLinkageAttr::CreateImplicit(AST));
5259 HLSLAssociatedResourceDeclAttr::CreateImplicit(AST, ResDecl));
5266static void handleArrayOfStructWithResources(
5268 EmbeddedResourceNameBuilder &NameBuilder, StructBindingContext &BindingCtx);
5273static void handleStructWithResources(
Sema &S,
VarDecl *ParentVD,
5275 EmbeddedResourceNameBuilder &NameBuilder,
5276 StructBindingContext &BindingCtx) {
5279 assert(RD->
getNumBases() <= 1 &&
"HLSL doesn't support multiple inheritance");
5286 handleStructWithResources(S, ParentVD, BaseRD, NameBuilder, BindingCtx);
5300 createGlobalResourceDeclForStruct(S, ParentVD, FD->
getLocation(), II,
5303 handleStructWithResources(S, ParentVD, RD, NameBuilder, BindingCtx);
5305 }
else if (
const auto *ArrayTy = dyn_cast<ConstantArrayType>(FDTy)) {
5307 "resource arrays should have been already handled");
5308 handleArrayOfStructWithResources(S, ParentVD, ArrayTy, NameBuilder,
5317handleArrayOfStructWithResources(
Sema &S,
VarDecl *ParentVD,
5319 EmbeddedResourceNameBuilder &NameBuilder,
5320 StructBindingContext &BindingCtx) {
5328 if (!SubCAT && !ElementRD)
5331 for (
unsigned I = 0, E = CAT->
getSize().getZExtValue(); I < E; ++I) {
5334 handleStructWithResources(S, ParentVD, ElementRD, NameBuilder,
5337 handleArrayOfStructWithResources(S, ParentVD, SubCAT, NameBuilder,
5350void SemaHLSL::handleGlobalStructOrArrayOfWithResources(
VarDecl *VD) {
5351 EmbeddedResourceNameBuilder NameBuilder(VD->
getName());
5352 StructBindingContext BindingCtx(VD);
5356 "Expected non-resource struct or array type");
5359 handleStructWithResources(
SemaRef, VD, RD, NameBuilder, BindingCtx);
5363 if (
const auto *CAT = dyn_cast<ConstantArrayType>(VDTy)) {
5364 handleArrayOfStructWithResources(
SemaRef, VD, CAT, NameBuilder, BindingCtx);
5372 if (
SemaRef.RequireCompleteType(
5375 diag::err_typecheck_decl_incomplete_type)) {
5389 DefaultCBufferDecls.push_back(VD);
5394 collectResourceBindingsOnVarDecl(VD);
5396 if (VD->
hasAttr<HLSLVkConstantIdAttr>())
5408 processExplicitBindingsOnDecl(VD);
5446 handleGlobalStructOrArrayOfWithResources(VD);
5450 if (VD->
hasAttr<HLSLGroupSharedAddressSpaceAttr>())
5459 "expected resource record type");
5475 const char *CreateMethodName;
5477 CreateMethodName = HasCounter ?
"__createFromBindingWithImplicitCounter"
5478 :
"__createFromBinding";
5480 CreateMethodName = HasCounter
5481 ?
"__createFromImplicitBindingWithImplicitCounter"
5482 :
"__createFromImplicitBinding";
5487 if (!CreateMethod) {
5492 "create method lookup should always succeed for built-in resource "
5501 Args.push_back(RegSlot);
5509 Args.push_back(OrderId);
5515 Args.push_back(Space);
5519 Args.push_back(RangeSize);
5523 Args.push_back(Index);
5525 StringRef VarName = VD->
getName();
5533 Args.push_back(NameCast);
5541 Args.push_back(CounterId);
5564 SemaRef.CheckCompleteVariableDeclaration(VD);
5570 "expected array of resource records");
5591 lookupMethod(
SemaRef, ResourceDecl,
5592 HasCounter ?
"__createFromBindingWithImplicitCounter"
5593 :
"__createFromBinding",
5597 CreateMethod = lookupMethod(
5599 HasCounter ?
"__createFromImplicitBindingWithImplicitCounter"
5600 :
"__createFromImplicitBinding",
5632std::optional<const DeclBindingInfo *> SemaHLSL::inferGlobalBinding(
Expr *E) {
5633 if (
auto *Ternary = dyn_cast<ConditionalOperator>(E)) {
5634 auto TrueInfo = inferGlobalBinding(Ternary->getTrueExpr());
5635 auto FalseInfo = inferGlobalBinding(Ternary->getFalseExpr());
5636 if (!TrueInfo || !FalseInfo)
5637 return std::nullopt;
5638 if (*TrueInfo != *FalseInfo)
5639 return std::nullopt;
5643 if (
auto *ASE = dyn_cast<ArraySubscriptExpr>(E))
5652 if (
const auto *AttrResType =
5653 HLSLAttributedResourceType::findHandleTypeOnResource(Ty)) {
5655 return Bindings.getDeclBindingInfo(VD, RC);
5662void SemaHLSL::trackLocalResource(
VarDecl *VD,
Expr *E) {
5663 std::optional<const DeclBindingInfo *> ExprBinding = inferGlobalBinding(E);
5666 diag::warn_hlsl_assigning_local_resource_is_not_unique)
5671 if (*ExprBinding ==
nullptr)
5674 auto PrevBinding = Assigns.find(VD);
5675 if (PrevBinding == Assigns.end()) {
5677 Assigns.insert({VD, *ExprBinding});
5682 if (*ExprBinding != PrevBinding->second) {
5684 diag::warn_hlsl_assigning_local_resource_is_not_unique)
5686 SemaRef.Diag(VD->getLocation(), diag::note_var_declared_here) << VD;
5697 "expected LHS to be a resource record or array of resource records");
5698 if (Opc != BO_Assign)
5703 while (
auto *ASE = dyn_cast<ArraySubscriptExpr>(E))
5711 SemaRef.Diag(Loc, diag::err_hlsl_assign_to_global_resource) << VD;
5716 trackLocalResource(VD, RHSExpr);
5724void SemaHLSL::collectResourceBindingsOnVarDecl(
VarDecl *VD) {
5726 "expected global variable that contains HLSL resource");
5729 if (
const HLSLBufferDecl *CBufferOrTBuffer = dyn_cast<HLSLBufferDecl>(VD)) {
5730 Bindings.addDeclBindingInfo(VD, CBufferOrTBuffer->isCBuffer()
5731 ? ResourceClass::CBuffer
5732 : ResourceClass::SRV);
5745 if (
const HLSLAttributedResourceType *AttrResType =
5746 HLSLAttributedResourceType::findHandleTypeOnResource(Ty)) {
5747 Bindings.addDeclBindingInfo(VD, AttrResType->getAttrs().ResourceClass);
5752 if (
const RecordType *RT = dyn_cast<RecordType>(Ty))
5753 collectResourceBindingsOnUserRecordDecl(VD, RT);
5759void SemaHLSL::processExplicitBindingsOnDecl(
VarDecl *VD) {
5762 bool HasBinding =
false;
5763 for (Attr *A : VD->
attrs()) {
5766 if (
auto PA = VD->
getAttr<HLSLVkPushConstantAttr>())
5767 Diag(PA->getLoc(), diag::err_hlsl_attr_incompatible) << A << PA;
5770 HLSLResourceBindingAttr *RBA = dyn_cast<HLSLResourceBindingAttr>(A);
5771 if (!RBA || !RBA->hasRegisterSlot())
5776 assert(RT != RegisterType::I &&
"invalid or obsolete register type should "
5777 "never have an attribute created");
5779 if (RT == RegisterType::C) {
5780 if (Bindings.hasBindingInfoForDecl(VD))
5782 diag::warn_hlsl_user_defined_type_missing_member)
5783 <<
static_cast<int>(RT);
5791 if (DeclBindingInfo *BI = Bindings.getDeclBindingInfo(VD, RC)) {
5796 diag::warn_hlsl_user_defined_type_missing_member)
5797 <<
static_cast<int>(RT);
5805class InitListTransformer {
5809 QualType *DstIt =
nullptr;
5810 Expr **ArgIt =
nullptr;
5816 bool castInitializer(Expr *E) {
5817 assert(DstIt &&
"This should always be something!");
5818 if (DstIt == DestTypes.end()) {
5820 ArgExprs.push_back(E);
5825 DstIt = DestTypes.begin();
5828 Ctx, *DstIt,
false);
5833 ArgExprs.push_back(
Init);
5838 bool buildInitializerListImpl(Expr *E) {
5840 if (
auto *
Init = dyn_cast<InitListExpr>(E)) {
5841 for (
auto *SubInit :
Init->inits())
5842 if (!buildInitializerListImpl(SubInit))
5852 return castInitializer(E);
5866 if (
auto *VecTy = Ty->
getAs<VectorType>()) {
5871 for (uint64_t I = 0; I <
Size; ++I) {
5873 SizeTy, SourceLocation());
5879 if (!castInitializer(ElExpr.
get()))
5884 if (
auto *MTy = Ty->
getAs<ConstantMatrixType>()) {
5885 unsigned Rows = MTy->getNumRows();
5886 unsigned Cols = MTy->getNumColumns();
5887 QualType ElemTy = MTy->getElementType();
5889 for (
unsigned R = 0;
R < Rows; ++
R) {
5890 for (
unsigned C = 0;
C < Cols; ++
C) {
5903 if (!castInitializer(ElExpr.
get()))
5911 if (
auto *ArrTy = dyn_cast<ConstantArrayType>(Ty.
getTypePtr())) {
5915 for (uint64_t I = 0; I <
Size; ++I) {
5917 SizeTy, SourceLocation());
5922 if (!buildInitializerListImpl(ElExpr.
get()))
5929 llvm::SmallVector<CXXRecordDecl *> RecordDecls;
5930 RecordDecls.push_back(RD);
5931 while (RecordDecls.back()->getNumBases()) {
5932 CXXRecordDecl *D = RecordDecls.back();
5934 "HLSL doesn't support multiple inheritance");
5935 RecordDecls.push_back(
5938 while (!RecordDecls.empty()) {
5939 CXXRecordDecl *RD = RecordDecls.pop_back_val();
5940 for (
auto *FD : RD->
fields()) {
5941 if (FD->isUnnamedBitField())
5949 if (!buildInitializerListImpl(Res.
get()))
5957 Expr *generateInitListsImpl(QualType Ty) {
5959 assert(ArgIt != ArgExprs.end() &&
"Something is off in iteration!");
5964 llvm::SmallVector<Expr *>
Inits;
5969 if (
auto *ATy = Ty->
getAs<VectorType>()) {
5970 ElTy = ATy->getElementType();
5971 Size = ATy->getNumElements();
5972 }
else if (
auto *CMTy = Ty->
getAs<ConstantMatrixType>()) {
5973 ElTy = CMTy->getElementType();
5974 Size = CMTy->getNumElementsFlattened();
5977 ElTy = VTy->getElementType();
5978 Size = VTy->getZExtSize();
5980 for (uint64_t I = 0; I <
Size; ++I)
5981 Inits.push_back(generateInitListsImpl(ElTy));
5984 llvm::SmallVector<CXXRecordDecl *> RecordDecls;
5985 RecordDecls.push_back(RD);
5986 while (RecordDecls.back()->getNumBases()) {
5987 CXXRecordDecl *D = RecordDecls.back();
5989 "HLSL doesn't support multiple inheritance");
5990 RecordDecls.push_back(
5993 while (!RecordDecls.empty()) {
5994 CXXRecordDecl *RD = RecordDecls.pop_back_val();
5995 for (
auto *FD : RD->
fields())
5996 if (!FD->isUnnamedBitField())
6000 auto *NewInit =
new (Ctx) InitListExpr(Ctx,
Inits.front()->getBeginLoc(),
6002 NewInit->setType(Ty);
6007 llvm::SmallVector<QualType, 16> DestTypes;
6008 llvm::SmallVector<Expr *, 16> ArgExprs;
6009 InitListTransformer(Sema &SemaRef,
const InitializedEntity &Entity)
6010 : S(SemaRef), Ctx(SemaRef.getASTContext()),
6011 Wrap(Entity.
getType()->isIncompleteArrayType()) {
6012 InitTy = Entity.
getType().getNonReferenceType();
6022 DstIt = DestTypes.begin();
6025 bool buildInitializerList(Expr *E) {
return buildInitializerListImpl(E); }
6027 Expr *generateInitLists() {
6028 assert(!ArgExprs.empty() &&
6029 "Call buildInitializerList to generate argument expressions.");
6030 ArgIt = ArgExprs.begin();
6032 return generateInitListsImpl(InitTy);
6033 llvm::SmallVector<Expr *>
Inits;
6034 while (ArgIt != ArgExprs.end())
6035 Inits.push_back(generateInitListsImpl(InitTy));
6037 auto *NewInit =
new (Ctx) InitListExpr(Ctx,
Inits.front()->getBeginLoc(),
6039 llvm::APInt ArySize(64,
Inits.size());
6041 ArraySizeModifier::Normal, 0));
6053 if (
const ArrayType *AT = dyn_cast<ArrayType>(Ty)) {
6060 if (
const auto *RT = Ty->
getAs<RecordType>()) {
6064 if (
const auto *CXXRD = dyn_cast<CXXRecordDecl>(RD)) {
6084 if (
Init->getType()->isScalarType())
6087 InitListTransformer ILT(
SemaRef, Entity);
6089 for (
unsigned I = 0; I <
Init->getNumInits(); ++I) {
6097 Init->setInit(I, E);
6099 if (!ILT.buildInitializerList(E))
6102 size_t ExpectedSize = ILT.DestTypes.size();
6103 size_t ActualSize = ILT.ArgExprs.size();
6104 if (ExpectedSize == 0 && ActualSize == 0)
6111 InitTy =
SemaRef.getASTContext().removeAddrSpaceQualType(InitTy);
6113 SemaRef.Diag(
Init->getBeginLoc(), diag::err_hlsl_incorrect_num_initializers)
6114 << (int)(ExpectedSize < ActualSize) << InitTy
6115 << ExpectedSize << ActualSize;
6125 assert(ExpectedSize > 0 &&
6126 "The expected size of an incomplete array type must be at least 1.");
6128 ((ActualSize + ExpectedSize - 1) / ExpectedSize) * ExpectedSize;
6136 InitTy =
SemaRef.getASTContext().removeAddrSpaceQualType(InitTy);
6137 if (ExpectedSize != ActualSize) {
6138 int TooManyOrFew = ActualSize > ExpectedSize ? 1 : 0;
6139 SemaRef.Diag(
Init->getBeginLoc(), diag::err_hlsl_incorrect_num_initializers)
6140 << TooManyOrFew << InitTy << ExpectedSize << ActualSize;
6147 Init->resizeInits(Ctx, NewInit->getNumInits());
6148 for (
unsigned I = 0; I < NewInit->getNumInits(); ++I)
6149 Init->updateInit(Ctx, I, NewInit->getInit(I));
6157 S.
Diag(OpLoc, diag::err_builtin_matrix_invalid_member)
6167 StringRef AccessorName = CompName->
getName();
6168 assert(!AccessorName.empty() &&
"Matrix Accessor must have a name");
6170 unsigned Rows = MT->getNumRows();
6171 unsigned Cols = MT->getNumColumns();
6172 bool IsZeroBasedAccessor =
false;
6173 unsigned ChunkLen = 0;
6174 if (AccessorName.size() < 2)
6176 "length 4 for zero based: \'_mRC\' or "
6177 "length 3 for one-based: \'_RC\' accessor",
6180 if (AccessorName[0] ==
'_') {
6181 if (AccessorName[1] ==
'm') {
6182 IsZeroBasedAccessor =
true;
6189 S, AccessorName,
"zero based: \'_mRC\' or one-based: \'_RC\' accessor",
6192 if (AccessorName.size() % ChunkLen != 0) {
6193 const llvm::StringRef
Expected = IsZeroBasedAccessor
6194 ?
"zero based: '_mRC' accessor"
6195 :
"one-based: '_RC' accessor";
6200 auto isDigit = [](
char c) {
return c >=
'0' &&
c <=
'9'; };
6201 auto isZeroBasedIndex = [](
unsigned i) {
return i <= 3; };
6202 auto isOneBasedIndex = [](
unsigned i) {
return i >= 1 && i <= 4; };
6204 bool HasRepeated =
false;
6206 unsigned NumComponents = 0;
6207 const char *Begin = AccessorName.data();
6209 for (
unsigned I = 0, E = AccessorName.size(); I < E; I += ChunkLen) {
6210 const char *Chunk = Begin + I;
6211 char RowChar = 0, ColChar = 0;
6212 if (IsZeroBasedAccessor) {
6214 if (Chunk[0] !=
'_' || Chunk[1] !=
'm') {
6215 char Bad = (Chunk[0] !=
'_') ? Chunk[0] : Chunk[1];
6217 S, StringRef(&Bad, 1),
"\'_m\' prefix",
6224 if (Chunk[0] !=
'_')
6226 S, StringRef(&Chunk[0], 1),
"\'_\' prefix",
6233 bool IsDigitsError =
false;
6235 unsigned BadPos = IsZeroBasedAccessor ? 2 : 1;
6239 IsDigitsError =
true;
6243 unsigned BadPos = IsZeroBasedAccessor ? 3 : 2;
6247 IsDigitsError =
true;
6252 unsigned Row = RowChar -
'0';
6253 unsigned Col = ColChar -
'0';
6255 bool HasIndexingError =
false;
6256 if (IsZeroBasedAccessor) {
6258 if (!isZeroBasedIndex(Row)) {
6259 S.
Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
6261 HasIndexingError =
true;
6263 if (!isZeroBasedIndex(Col)) {
6264 S.
Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
6266 HasIndexingError =
true;
6270 if (!isOneBasedIndex(Row)) {
6271 S.
Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
6273 HasIndexingError =
true;
6275 if (!isOneBasedIndex(Col)) {
6276 S.
Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
6278 HasIndexingError =
true;
6285 if (HasIndexingError)
6291 bool HasBoundsError =
false;
6293 Diag(OpLoc, diag::err_hlsl_matrix_index_out_of_bounds)
6295 HasBoundsError =
true;
6298 Diag(OpLoc, diag::err_hlsl_matrix_index_out_of_bounds)
6300 HasBoundsError =
true;
6305 unsigned FlatIndex = Row * Cols + Col;
6306 if (Seen[FlatIndex])
6308 Seen[FlatIndex] =
true;
6311 if (NumComponents == 0 || NumComponents > 4) {
6312 S.
Diag(OpLoc, diag::err_hlsl_matrix_swizzle_invalid_length)
6317 QualType ElemTy = MT->getElementType();
6318 if (NumComponents == 1)
6324 for (Sema::ExtVectorDeclsType::iterator
6328 if ((*I)->getUnderlyingType() == VT)
6339 trackLocalResource(VDecl,
Init);
6341 const HLSLVkConstantIdAttr *ConstIdAttr =
6342 VDecl->
getAttr<HLSLVkConstantIdAttr>();
6349 if (!
Init->isCXX11ConstantExpr(Context, &InitValue)) {
6359 int ConstantID = ConstIdAttr->getId();
6360 llvm::APInt IDVal(Context.getIntWidth(Context.IntTy), ConstantID);
6362 ConstIdAttr->getLocation());
6366 if (
C->getType()->getCanonicalTypeUnqualified() !=
6370 Context.getTrivialTypeSourceInfo(
6371 Init->getType(),
Init->getExprLoc()),
6390 if (!Params || Params->
size() != 1)
6403 if (
auto *TTP = dyn_cast<TemplateTypeParmDecl>(P)) {
6404 if (TTP->hasDefaultArgument()) {
6405 TemplateArgs.
addArgument(TTP->getDefaultArgument());
6408 }
else if (
auto *NTTP = dyn_cast<NonTypeTemplateParmDecl>(P)) {
6409 if (NTTP->hasDefaultArgument()) {
6410 TemplateArgs.
addArgument(NTTP->getDefaultArgument());
6413 }
else if (
auto *TTPD = dyn_cast<TemplateTemplateParmDecl>(P)) {
6414 if (TTPD->hasDefaultArgument()) {
6415 TemplateArgs.
addArgument(TTPD->getDefaultArgument());
6422 return SemaRef.CheckTemplateIdType(
6424 TemplateArgs,
nullptr,
false);
Defines the clang::ASTContext interface.
Defines enum values for all the target-independent builtin functions.
llvm::dxil::ResourceClass ResourceClass
Defines the C++ Decl subclasses, other than those for templates (found in DeclTemplate....
Defines the clang::IdentifierInfo, clang::IdentifierTable, and clang::Selector interfaces.
Forward-declares and imports various common LLVM datatypes that clang wants to use unqualified.
llvm::SmallVector< std::pair< const MemRegion *, SVal >, 4 > Bindings
static bool CheckArgTypeMatches(Sema *S, Expr *Arg, QualType ExpectedType)
static void BuildFlattenedTypeList(QualType BaseTy, llvm::SmallVectorImpl< QualType > &List)
static bool CheckUnsignedIntRepresentation(Sema *S, SourceLocation Loc, int ArgOrdinal, clang::QualType PassedType)
static bool containsIncompleteArrayType(QualType Ty)
static QualType handleIntegerVectorBinOpConversion(Sema &SemaRef, ExprResult &LHS, ExprResult &RHS, QualType LHSType, QualType RHSType, QualType LElTy, QualType RElTy, bool IsCompAssign)
static bool convertToRegisterType(StringRef Slot, RegisterType *RT)
static StringRef createRegisterString(ASTContext &AST, RegisterType RegType, unsigned N)
static bool CheckWaveActive(Sema *S, CallExpr *TheCall)
static void createHostLayoutStructForBuffer(Sema &S, HLSLBufferDecl *BufDecl)
static void castVector(Sema &S, ExprResult &E, QualType &Ty, unsigned Sz)
static QualType ReportMatrixInvalidMember(Sema &S, StringRef Name, StringRef Expected, SourceLocation OpLoc, SourceLocation CompLoc)
static bool CheckBoolSelect(Sema *S, CallExpr *TheCall)
static unsigned calculateLegacyCbufferFieldAlign(const ASTContext &Context, QualType T)
static bool isZeroSizedArray(const ConstantArrayType *CAT)
static bool DiagnoseHLSLRegisterAttribute(Sema &S, SourceLocation &ArgLoc, Decl *D, RegisterType RegType, bool SpecifiedSpace)
static bool hasConstantBufferLayout(QualType QT)
static FieldDecl * createFieldForHostLayoutStruct(Sema &S, const Type *Ty, IdentifierInfo *II, CXXRecordDecl *LayoutStruct)
static bool CheckUnsignedIntVecRepresentation(Sema *S, SourceLocation Loc, int ArgOrdinal, clang::QualType PassedType)
static bool isInvalidConstantBufferLeafElementType(const Type *Ty)
static bool CheckCalculateLodBuiltin(Sema &S, CallExpr *TheCall)
static Builtin::ID getSpecConstBuiltinId(const Type *Type)
static bool CheckFloatingOrIntRepresentation(Sema *S, SourceLocation Loc, int ArgOrdinal, clang::QualType PassedType)
static const Type * createHostLayoutType(Sema &S, const Type *Ty)
static bool CheckAnyScalarOrVector(Sema *S, CallExpr *TheCall, unsigned ArgIndex)
static const HLSLAttributedResourceType * getResourceArrayHandleType(QualType QT)
static IdentifierInfo * getHostLayoutStructName(Sema &S, NamedDecl *BaseDecl, bool MustBeUnique)
static void addImplicitBindingAttrToDecl(Sema &S, Decl *D, RegisterType RT, uint32_t ImplicitBindingOrderID)
static void SetElementTypeAsReturnType(Sema *S, CallExpr *TheCall, QualType ReturnType)
static unsigned calculateLegacyCbufferSize(const ASTContext &Context, QualType T)
static bool CheckLoadLevelBuiltin(Sema &S, CallExpr *TheCall)
static RegisterType getRegisterType(ResourceClass RC)
static bool ValidateRegisterNumber(uint64_t SlotNum, Decl *TheDecl, ASTContext &Ctx, RegisterType RegTy)
static bool isVkPipelineBuiltin(const ASTContext &AstContext, FunctionDecl *FD, HLSLAppliedSemanticAttr *Semantic, bool IsInput)
static bool CheckVectorElementCount(Sema *S, QualType PassedType, QualType BaseType, unsigned ExpectedCount, SourceLocation Loc)
static bool CheckModifiableLValue(Sema *S, CallExpr *TheCall, unsigned ArgIndex)
static QualType castElement(Sema &S, ExprResult &E, QualType Ty)
static char getRegisterTypeChar(RegisterType RT)
static bool CheckNotBoolScalarOrVector(Sema *S, CallExpr *TheCall, unsigned ArgIndex)
static bool isMatrixOrArrayOfMatrix(const ASTContext &Ctx, QualType QT)
static CXXRecordDecl * findRecordDeclInContext(IdentifierInfo *II, DeclContext *DC)
static bool CheckWavePrefix(Sema *S, CallExpr *TheCall)
static bool CheckExpectedBitWidth(Sema *S, CallExpr *TheCall, unsigned ArgOrdinal, unsigned Width)
static LangAS getLangASFromResourceClass(ResourceClass RC)
static bool diagnoseMatrixLayoutOnNonMatrix(Sema &SemaRef, Decl *D, SourceLocation Loc, const IdentifierInfo *AttrName)
static bool CheckVectorSelect(Sema *S, CallExpr *TheCall)
static QualType handleFloatVectorBinOpConversion(Sema &SemaRef, ExprResult &LHS, ExprResult &RHS, QualType LHSType, QualType RHSType, QualType LElTy, QualType RElTy, bool IsCompAssign)
static ResourceClass getResourceClass(RegisterType RT)
static CXXRecordDecl * createHostLayoutStruct(Sema &S, CXXRecordDecl *StructDecl)
static bool CheckScalarOrVector(Sema *S, CallExpr *TheCall, QualType Scalar, unsigned ArgIndex)
static bool CheckSamplingBuiltin(Sema &S, CallExpr *TheCall, SampleKind Kind)
static bool CheckScalarOrVectorOrMatrix(Sema *S, CallExpr *TheCall, QualType Scalar, unsigned ArgIndex)
static bool CheckFloatRepresentation(Sema *S, SourceLocation Loc, int ArgOrdinal, clang::QualType PassedType)
static bool CheckAnyDoubleRepresentation(Sema *S, SourceLocation Loc, int ArgOrdinal, clang::QualType PassedType)
static bool requiresImplicitBufferLayoutStructure(const CXXRecordDecl *RD)
static bool CheckResourceHandle(Sema *S, CallExpr *TheCall, unsigned ArgIndex, llvm::function_ref< bool(const HLSLAttributedResourceType *ResType)> Check=nullptr)
static void validatePackoffset(Sema &S, HLSLBufferDecl *BufDecl)
static bool IsDefaultBufferConstantDecl(const ASTContext &Ctx, VarDecl *VD)
HLSLResourceBindingAttr::RegisterType RegisterType
static CastKind getScalarCastKind(ASTContext &Ctx, QualType DestTy, QualType SrcTy)
static bool CheckGatherBuiltin(Sema &S, CallExpr *TheCall, bool IsCmp)
static bool isValidWaveSizeValue(unsigned Value)
static bool isResourceRecordTypeOrArrayOf(QualType Ty)
static bool AccumulateHLSLResourceSlots(QualType Ty, uint64_t &StartSlot, const uint64_t &Limit, const ResourceClass ResClass, ASTContext &Ctx, uint64_t ArrayCount=1)
static bool CheckNoDoubleVectors(Sema *S, SourceLocation Loc, int ArgOrdinal, clang::QualType PassedType)
static bool ValidateMultipleRegisterAnnotations(Sema &S, Decl *TheDecl, RegisterType regType)
static bool CheckTextureSamplerAndLocation(Sema &S, CallExpr *TheCall)
static bool DiagnoseLocalRegisterBinding(Sema &S, SourceLocation &ArgLoc, Decl *D, RegisterType RegType, bool SpecifiedSpace)
static bool CheckIndexType(Sema *S, CallExpr *TheCall, unsigned IndexArgIndex)
This file declares semantic analysis for HLSL constructs.
Defines the clang::SourceLocation class and associated facilities.
Defines various enumerations that describe declaration and type specifiers.
C Language Family Type Representation.
Defines the clang::TypeLoc interface and its subclasses.
C Language Family Type Representation.
static const TypeInfo & getInfo(unsigned id)
__device__ __2f16 float c
return(__x > > __y)|(__x<<(32 - __y))
APValue - This class implements a discriminated union of [uninitialized] [APSInt] [APFloat],...
virtual bool HandleTopLevelDecl(DeclGroupRef D)
HandleTopLevelDecl - Handle the specified top-level declaration.
Holds long-lived AST nodes (such as types and decls) that can be referred to throughout the semantic ...
unsigned getIntWidth(QualType T) const
int getIntegerTypeOrder(QualType LHS, QualType RHS) const
Return the highest ranked integer type, see C99 6.3.1.8p1.
QualType getPointerType(QualType T) const
Return the uniqued reference to the type for a pointer to the specified type.
const IncompleteArrayType * getAsIncompleteArrayType(QualType T) const
QualType getConstantArrayType(QualType EltTy, const llvm::APInt &ArySize, const Expr *SizeExpr, ArraySizeModifier ASM, unsigned IndexTypeQuals) const
Return the unique reference to the type for a constant array of the specified element type.
QualType getBaseElementType(const ArrayType *VAT) const
Return the innermost element type of an array type.
int getFloatingTypeOrder(QualType LHS, QualType RHS) const
Compare the rank of the two specified floating point types, ignoring the domain of the type (i....
TypeSourceInfo * getTrivialTypeSourceInfo(QualType T, SourceLocation Loc=SourceLocation()) const
Allocate a TypeSourceInfo where all locations have been initialized to a given location,...
QualType getStringLiteralArrayType(QualType EltTy, unsigned Length) const
Return a type for a constant array for a string literal of the specified element type and length.
uint64_t getTypeSize(QualType T) const
Return the size of the specified (complete) type T, in bits.
CharUnits getTypeSizeInChars(QualType T) const
Return the size of the specified (complete) type T, in characters.
CanQualType UnsignedIntTy
QualType getTypedefType(ElaboratedTypeKeyword Keyword, NestedNameSpecifier Qualifier, const TypedefNameDecl *Decl, QualType UnderlyingType=QualType(), std::optional< bool > TypeMatchesDeclOrNone=std::nullopt) const
Return the unique reference to the type for the specified typedef-name decl.
llvm::StringRef backupStr(llvm::StringRef S) const
QualType getSizeType() const
Return the unique type for "size_t" (C99 7.17), defined in <stddef.h>.
QualType getExtVectorType(QualType VectorType, unsigned NumElts) const
Return the unique reference to an extended vector type of the specified element type and size.
const TargetInfo & getTargetInfo() const
QualType getHLSLAttributedResourceType(QualType Wrapped, QualType Contained, const HLSLAttributedResourceType::Attributes &Attrs)
QualType getAddrSpaceQualType(QualType T, LangAS AddressSpace) const
Return the uniqued reference to the type for an address space qualified type with the specified type ...
CanQualType getCanonicalTagType(const TagDecl *TD) const
static bool hasSameUnqualifiedType(QualType T1, QualType T2)
Determine whether the given types are equivalent after cvr-qualifiers have been removed.
QualType getConstantMatrixType(QualType ElementType, unsigned NumRows, unsigned NumColumns) const
Return the unique reference to the matrix type of the specified element type and size.
unsigned getTypeAlign(QualType T) const
Return the ABI-specified alignment of a (complete) type T, in bits.
Represents an array type, per C99 6.7.5.2 - Array Declarators.
QualType getElementType() const
Attr - This represents one attribute.
attr::Kind getKind() const
SourceLocation getLocation() const
SourceLocation getScopeLoc() const
SourceRange getRange() const
const IdentifierInfo * getScopeName() const
SourceLocation getLoc() const
const IdentifierInfo * getAttrName() const
Represents a base class of a C++ class.
QualType getType() const
Retrieves the type of the base class.
Represents a static or instance method of a struct/union/class.
Represents a C++ struct/union/class.
bool isHLSLIntangible() const
Returns true if the class contains HLSL intangible type, either as a field or in base class.
static CXXRecordDecl * Create(const ASTContext &C, TagKind TK, DeclContext *DC, SourceLocation StartLoc, SourceLocation IdLoc, IdentifierInfo *Id, CXXRecordDecl *PrevDecl=nullptr)
void setBases(CXXBaseSpecifier const *const *Bases, unsigned NumBases)
Sets the base classes of this struct or class.
base_class_iterator bases_end()
void completeDefinition() override
Indicates that the definition of this class is now complete.
unsigned getNumBases() const
Retrieves the number of base classes of this class.
base_class_iterator bases_begin()
bool isEmpty() const
Determine whether this is an empty class in the sense of (C++11 [meta.unary.prop]).
CallExpr - Represents a function call (C99 6.5.2.2, C++ [expr.call]).
Expr * getArg(unsigned Arg)
getArg - Return the specified argument.
SourceLocation getBeginLoc() const
static CallExpr * Create(const ASTContext &Ctx, Expr *Fn, ArrayRef< Expr * > Args, QualType Ty, ExprValueKind VK, SourceLocation RParenLoc, FPOptionsOverride FPFeatures, unsigned MinNumArgs=0, ADLCallKind UsesADL=NotADL)
Create a call expression.
FunctionDecl * getDirectCallee()
If the callee is a FunctionDecl, return it. Otherwise return null.
unsigned getNumArgs() const
getNumArgs - Return the number of actual arguments to this call.
static CanQual< Type > CreateUnsafe(QualType Other)
QualType withConst() const
Retrieves a version of this type with const applied.
const T * getTypePtr() const
Retrieve the underlying type pointer, which refers to a canonical type.
QuantityType getQuantity() const
getQuantity - Get the raw integer representation of this quantity.
Represents the canonical version of C arrays with a specified constant size.
bool isZeroSize() const
Return true if the size is zero.
llvm::APInt getSize() const
Return the constant array size as an APInt.
uint64_t getZExtSize() const
Return the size zero-extended as a uint64_t.
Represents a concrete matrix type with constant number of rows and columns.
unsigned getNumColumns() const
Returns the number of columns in the matrix.
static DeclAccessPair make(NamedDecl *D, AccessSpecifier AS)
DeclContext - This is used only as base class of specific decl types that can act as declaration cont...
lookup_result lookup(DeclarationName Name) const
lookup - Find the declarations (if any) with the given Name in this context.
bool isTranslationUnit() const
void addDecl(Decl *D)
Add the declaration D into this context.
decl_range decls() const
decls_begin/decls_end - Iterate over the declarations stored in this context.
DeclContext * getNonTransparentContext()
A reference to a declared variable, function, enum, etc.
static DeclRefExpr * Create(const ASTContext &Context, NestedNameSpecifierLoc QualifierLoc, SourceLocation TemplateKWLoc, ValueDecl *D, bool RefersToEnclosingVariableOrCapture, SourceLocation NameLoc, QualType T, ExprValueKind VK, NamedDecl *FoundD=nullptr, const TemplateArgumentListInfo *TemplateArgs=nullptr, NonOdrUseReason NOUR=NOUR_None)
Decl - This represents one declaration (or definition), e.g.
ASTContext & getASTContext() const LLVM_READONLY
attr_iterator attr_end() const
bool isImplicit() const
isImplicit - Indicates whether the declaration was implicitly generated by the implementation.
void setInvalidDecl(bool Invalid=true)
setInvalidDecl - Indicates the Decl had a semantic error.
bool isInExportDeclContext() const
Whether this declaration was exported in a lexical context.
attr_iterator attr_begin() const
DeclContext * getNonTransparentDeclContext()
Return the non transparent context.
SourceLocation getLocation() const
void setImplicit(bool I=true)
DeclContext * getDeclContext()
AccessSpecifier getAccess() const
SourceLocation getBeginLoc() const LLVM_READONLY
The name of a declaration.
Represents a ValueDecl that came out of a declarator.
SourceLocation getBeginLoc() const LLVM_READONLY
This represents one expression.
bool isIntegerConstantExpr(const ASTContext &Ctx) const
ExprValueKind getValueKind() const
getValueKind - The value kind that this expression produces.
Expr * IgnoreParenImpCasts() LLVM_READONLY
Skip past any parentheses and implicit casts which might surround this expression until reaching a fi...
Expr * IgnoreParens() LLVM_READONLY
Skip past any parentheses which might surround this expression until reaching a fixed point.
std::optional< llvm::APSInt > getIntegerConstantExpr(const ASTContext &Ctx) const
isIntegerConstantExpr - Return the value if this expression is a valid integer constant expression.
bool isLValue() const
isLValue - True if this expression is an "l-value" according to the rules of the current language.
ExprObjectKind getObjectKind() const
getObjectKind - The object kind that this expression produces.
bool HasSideEffects(const ASTContext &Ctx, bool IncludePossibleEffects=true) const
HasSideEffects - This routine returns true for all those expressions which have any effect other than...
SourceLocation getExprLoc() const LLVM_READONLY
getExprLoc - Return the preferred location for the arrow when diagnosing a problem with a generic exp...
ExtVectorType - Extended vector type.
Represents difference between two FPOptions values.
Represents a member of a struct/union/class.
static FieldDecl * Create(const ASTContext &C, DeclContext *DC, SourceLocation StartLoc, SourceLocation IdLoc, const IdentifierInfo *Id, QualType T, TypeSourceInfo *TInfo, Expr *BW, bool Mutable, InClassInitStyle InitStyle)
static FixItHint CreateReplacement(CharSourceRange RemoveRange, StringRef Code)
Create a code modification hint that replaces the given source range with the given code string.
Represents a function declaration or definition.
const ParmVarDecl * getParamDecl(unsigned i) const
Stmt * getBody(const FunctionDecl *&Definition) const
Retrieve the body (definition) of the function.
bool isThisDeclarationADefinition() const
Returns whether this specific declaration of the function is also a definition that does not contain ...
QualType getReturnType() const
ArrayRef< ParmVarDecl * > parameters() const
bool isTemplateInstantiation() const
Determines if the given function was instantiated from a function template.
redecl_range redecls() const
Returns an iterator range for all the redeclarations of the same decl.
unsigned getNumParams() const
Return the number of parameters this function must have based on its FunctionType.
DeclarationNameInfo getNameInfo() const
bool hasBody(const FunctionDecl *&Definition) const
Returns true if the function has a body.
bool isDefined(const FunctionDecl *&Definition, bool CheckForPendingFriendDefinition=false) const
Returns true if the function has a definition that does not need to be instantiated.
Represents a prototype with parameter type info, e.g.
HLSLBufferDecl - Represent a cbuffer or tbuffer declaration.
static HLSLBufferDecl * Create(ASTContext &C, DeclContext *LexicalParent, bool CBuffer, SourceLocation KwLoc, IdentifierInfo *ID, SourceLocation IDLoc, SourceLocation LBrace)
void addLayoutStruct(CXXRecordDecl *LS)
void setHasValidPackoffset(bool PO)
static HLSLBufferDecl * CreateDefaultCBuffer(ASTContext &C, DeclContext *LexicalParent, ArrayRef< Decl * > DefaultCBufferDecls)
buffer_decl_range buffer_decls() const
static HLSLOutArgExpr * Create(const ASTContext &C, QualType Ty, OpaqueValueExpr *Base, OpaqueValueExpr *OpV, Expr *WB, bool IsInOut)
static HLSLRootSignatureDecl * Create(ASTContext &C, DeclContext *DC, SourceLocation Loc, IdentifierInfo *ID, llvm::dxbc::RootSignatureVersion Version, ArrayRef< llvm::hlsl::rootsig::RootElement > RootElements)
One of these records is kept for each identifier that is lexed.
StringRef getName() const
Return the actual identifier string.
A simple pair of identifier info and location.
SourceLocation getLoc() const
IdentifierInfo * getIdentifierInfo() const
IdentifierInfo & get(StringRef Name)
Return the identifier token info for the specified named identifier.
ImplicitCastExpr - Allows us to explicitly represent implicit type conversions, which have no direct ...
static ImplicitCastExpr * Create(const ASTContext &Context, QualType T, CastKind Kind, Expr *Operand, const CXXCastPath *BasePath, ExprValueKind Cat, FPOptionsOverride FPO)
Describes an C or C++ initializer list.
Describes an entity that is being initialized.
QualType getType() const
Retrieve type being initialized.
static InitializedEntity InitializeParameter(ASTContext &Context, ParmVarDecl *Parm)
Create the initialization entity for a parameter.
static IntegerLiteral * Create(const ASTContext &C, const llvm::APInt &V, QualType type, SourceLocation l)
Returns a new integer literal with value 'V' and type 'type'.
iterator begin(Source *source, bool LocalOnly=false)
Represents the results of name lookup.
Represents a prvalue temporary that is written into memory so that a reference can bind to it.
Represents a matrix type, as defined in the Matrix Types clang extensions.
MemberExpr - [C99 6.5.2.3] Structure and Union Members.
ValueDecl * getMemberDecl() const
Retrieve the member declaration to which this expression refers.
This represents a decl that may have a name.
NamedDecl * getUnderlyingDecl()
Looks through UsingDecls and ObjCCompatibleAliasDecls for the underlying named decl.
IdentifierInfo * getIdentifier() const
Get the identifier that names this declaration, if there is one.
StringRef getName() const
Get the name of identifier for this declaration as a StringRef.
DeclarationName getDeclName() const
Get the actual, stored name of the declaration, which may be a special name.
A C++ nested-name-specifier augmented with source location information.
OpaqueValueExpr - An expression referring to an opaque object of a fixed type and value class.
Represents a parameter to a function.
ParsedAttr - Represents a syntactic attribute.
unsigned getSemanticSpelling() const
If the parsed attribute has a semantic equivalent, and it would have a semantic Spelling enumeration ...
unsigned getMinArgs() const
bool checkExactlyNumArgs(class Sema &S, unsigned Num) const
Check if the attribute has exactly as many args as Num.
IdentifierLoc * getArgAsIdent(unsigned Arg) const
bool hasParsedType() const
const ParsedType & getTypeArg() const
unsigned getNumArgs() const
getNumArgs - Return the number of actual arguments to this attribute.
bool isArgIdent(unsigned Arg) const
Expr * getArgAsExpr(unsigned Arg) const
AttributeCommonInfo::Kind getKind() const
A (possibly-)qualified type.
void addRestrict()
Add the restrict qualifier to this QualType.
QualType getNonLValueExprType(const ASTContext &Context) const
Determine the type of a (typically non-lvalue) expression with the specified result type.
QualType getDesugaredType(const ASTContext &Context) const
Return the specified type with any "sugar" removed from the type.
bool isNull() const
Return true if this QualType doesn't point to a type yet.
const Type * getTypePtr() const
Retrieves a pointer to the underlying (unqualified) type.
LangAS getAddressSpace() const
Return the address space of this type.
QualType getNonReferenceType() const
If Type is a reference type (e.g., const int&), returns the type that the reference refers to ("const...
QualType getCanonicalType() const
QualType getUnqualifiedType() const
Retrieve the unqualified variant of the given type, removing as little sugar as possible.
bool hasAddressSpace() const
Check if this type has any address space qualifier.
Represents a struct/union/class.
field_range fields() const
bool hasBindingInfoForDecl(const VarDecl *VD) const
DeclBindingInfo * getDeclBindingInfo(const VarDecl *VD, ResourceClass ResClass)
DeclBindingInfo * addDeclBindingInfo(const VarDecl *VD, ResourceClass ResClass)
Scope - A scope is a transient data structure that is used while parsing the program.
ASTContext & getASTContext() const
SemaDiagnosticBuilder Diag(SourceLocation Loc, unsigned DiagID)
Emit a diagnostic.
ExprResult ActOnOutParamExpr(ParmVarDecl *Param, Expr *Arg)
HLSLRootSignatureDecl * lookupRootSignatureOverrideDecl(DeclContext *DC) const
bool CanPerformElementwiseCast(Expr *Src, QualType DestType)
void handleWaveSizeAttr(Decl *D, const ParsedAttr &AL)
void handleVkLocationAttr(Decl *D, const ParsedAttr &AL)
HLSLAttributedResourceLocInfo TakeLocForHLSLAttribute(const HLSLAttributedResourceType *RT)
void handleSemanticAttr(Decl *D, const ParsedAttr &AL)
bool CanPerformScalarCast(QualType SrcTy, QualType DestTy)
QualType ProcessResourceTypeAttributes(QualType Wrapped)
void handleShaderAttr(Decl *D, const ParsedAttr &AL)
uint32_t getNextImplicitBindingOrderID()
void CheckEntryPoint(FunctionDecl *FD)
void handleVkExtBuiltinOutputAttr(Decl *D, const ParsedAttr &AL)
void emitLogicalOperatorFixIt(Expr *LHS, Expr *RHS, BinaryOperatorKind Opc)
T * createSemanticAttr(const AttributeCommonInfo &ACI, std::optional< unsigned > Location)
bool initGlobalResourceDecl(VarDecl *VD)
void ActOnEndOfTranslationUnit(TranslationUnitDecl *TU)
bool initGlobalResourceArrayDecl(VarDecl *VD)
HLSLVkConstantIdAttr * mergeVkConstantIdAttr(Decl *D, const AttributeCommonInfo &AL, int Id)
HLSLNumThreadsAttr * mergeNumThreadsAttr(Decl *D, const AttributeCommonInfo &AL, int X, int Y, int Z)
void deduceAddressSpace(VarDecl *Decl)
std::pair< IdentifierInfo *, bool > ActOnStartRootSignatureDecl(StringRef Signature)
Computes the unique Root Signature identifier from the given signature, then lookup if there is a pre...
void handlePackOffsetAttr(Decl *D, const ParsedAttr &AL)
std::optional< ExprResult > tryPerformConstantBufferConversion(ExprResult &BaseExpr)
bool diagnosePositionType(QualType T, const ParsedAttr &AL)
bool handleInitialization(VarDecl *VDecl, Expr *&Init)
bool diagnoseInputIDType(QualType T, const ParsedAttr &AL)
void handleParamModifierAttr(Decl *D, const ParsedAttr &AL)
bool CheckResourceBinOp(BinaryOperatorKind Opc, Expr *LHSExpr, Expr *RHSExpr, SourceLocation Loc)
bool CanPerformAggregateSplatCast(Expr *Src, QualType DestType)
bool ActOnResourceMemberAccessExpr(MemberExpr *ME)
bool IsScalarizedLayoutCompatible(QualType T1, QualType T2) const
QualType ActOnTemplateShorthand(TemplateDecl *Template, SourceLocation NameLoc)
void diagnoseSystemSemanticAttr(Decl *D, const ParsedAttr &AL, std::optional< unsigned > Index)
void handleRootSignatureAttr(Decl *D, const ParsedAttr &AL)
bool CheckCompatibleParameterABI(FunctionDecl *New, FunctionDecl *Old)
QualType handleVectorBinOpConversion(ExprResult &LHS, ExprResult &RHS, QualType LHSType, QualType RHSType, bool IsCompAssign)
QualType checkMatrixComponent(Sema &S, QualType baseType, ExprValueKind &VK, SourceLocation OpLoc, const IdentifierInfo *CompName, SourceLocation CompLoc)
void handleResourceBindingAttr(Decl *D, const ParsedAttr &AL)
bool IsTypedResourceElementCompatible(QualType T1)
bool transformInitList(const InitializedEntity &Entity, InitListExpr *Init)
void handleNumThreadsAttr(Decl *D, const ParsedAttr &AL)
bool ActOnUninitializedVarDecl(VarDecl *D)
void handleVkExtBuiltinInputAttr(Decl *D, const ParsedAttr &AL)
bool diagnoseInstantiatedMatrixLayoutAttr(Decl *D, const HLSLMatrixLayoutAttr *Attr)
void ActOnTopLevelFunction(FunctionDecl *FD)
bool handleResourceTypeAttr(QualType T, const ParsedAttr &AL)
void handleVkPushConstantAttr(Decl *D, const ParsedAttr &AL)
HLSLShaderAttr * mergeShaderAttr(Decl *D, const AttributeCommonInfo &AL, llvm::Triple::EnvironmentType ShaderType)
NamedDecl * getConstantBufferConversionFunction(QualType Type, CXXRecordDecl *RD)
void ActOnFinishBuffer(Decl *Dcl, SourceLocation RBrace)
void handleVkBindingAttr(Decl *D, const ParsedAttr &AL)
HLSLParamModifierAttr * mergeParamModifierAttr(Decl *D, const AttributeCommonInfo &AL, HLSLParamModifierAttr::Spelling Spelling)
QualType getInoutParameterType(QualType Ty)
void handleMatrixLayoutAttr(Decl *D, const ParsedAttr &AL)
void handleVkConstantIdAttr(Decl *D, const ParsedAttr &AL)
Decl * ActOnStartBuffer(Scope *BufferScope, bool CBuffer, SourceLocation KwLoc, IdentifierInfo *Ident, SourceLocation IdentLoc, SourceLocation LBrace)
HLSLWaveSizeAttr * mergeWaveSizeAttr(Decl *D, const AttributeCommonInfo &AL, int Min, int Max, int Preferred, int SpelledArgsCount)
bool handleRootSignatureElements(ArrayRef< hlsl::RootSignatureElement > Elements)
void ActOnFinishRootSignatureDecl(SourceLocation Loc, IdentifierInfo *DeclIdent, ArrayRef< hlsl::RootSignatureElement > Elements)
Creates the Root Signature decl of the parsed Root Signature elements onto the AST and push it onto c...
void ActOnVariableDeclarator(VarDecl *VD)
bool CheckBuiltinFunctionCall(unsigned BuiltinID, CallExpr *TheCall)
Sema - This implements semantic analysis and AST building for C.
@ LookupOrdinaryName
Ordinary name lookup, which finds ordinary names (functions, variables, typedefs, etc....
@ LookupMemberName
Member name lookup, which finds the names of class/struct/union members.
ExtVectorDeclsType ExtVectorDecls
ExtVectorDecls - This is a list all the extended vector types.
ASTContext & getASTContext() const
ExprResult ImpCastExprToType(Expr *E, QualType Type, CastKind CK, ExprValueKind VK=VK_PRValue, const CXXCastPath *BasePath=nullptr, CheckedConversionKind CCK=CheckedConversionKind::Implicit)
ImpCastExprToType - If Expr is not of type 'Type', insert an implicit cast.
const LangOptions & getLangOpts() const
ExprResult TemporaryMaterializationConversion(Expr *E)
If E is a prvalue denoting an unmaterialized temporary, materialize it as an xvalue.
ExprResult BuildFieldReferenceExpr(Expr *BaseExpr, bool IsArrow, SourceLocation OpLoc, const CXXScopeSpec &SS, FieldDecl *Field, DeclAccessPair FoundDecl, const DeclarationNameInfo &MemberNameInfo)
bool checkArgCountRange(CallExpr *Call, unsigned MinArgCount, unsigned MaxArgCount)
Checks that a call expression's argument count is in the desired range.
ExternalSemaSource * getExternalSource() const
bool checkArgCount(CallExpr *Call, unsigned DesiredArgCount)
Checks that a call expression's argument count is the desired number.
ExprResult CreateBuiltinArraySubscriptExpr(Expr *Base, SourceLocation LLoc, Expr *Idx, SourceLocation RLoc)
bool LookupQualifiedName(LookupResult &R, DeclContext *LookupCtx, bool InUnqualifiedLookup=false)
Perform qualified name lookup into a given context.
ExprResult PerformCopyInitialization(const InitializedEntity &Entity, SourceLocation EqualLoc, ExprResult Init, bool TopLevelOfInitList=false, bool AllowExplicit=false)
ExprResult CreateBuiltinMatrixSubscriptExpr(Expr *Base, Expr *RowIdx, Expr *ColumnIdx, SourceLocation RBLoc)
Encodes a location in the source.
SourceLocation getLocWithOffset(IntTy Offset) const
Return a source location with the specified offset from this SourceLocation.
A trivial tuple used to represent a source range.
SourceLocation getEnd() const
SourceLocation getEndLoc() const LLVM_READONLY
void printPretty(raw_ostream &OS, PrinterHelper *Helper, const PrintingPolicy &Policy, unsigned Indentation=0, StringRef NewlineSymbol="\n", const ASTContext *Context=nullptr) const
SourceRange getSourceRange() const LLVM_READONLY
SourceLocation tokens are not useful in isolation - they are low level value objects created/interpre...
SourceLocation getBeginLoc() const LLVM_READONLY
StringLiteral - This represents a string literal expression, e.g.
static StringLiteral * Create(const ASTContext &Ctx, StringRef Str, StringLiteralKind Kind, bool Pascal, QualType Ty, ArrayRef< SourceLocation > Locs)
This is the "fully general" constructor that allows representation of strings formed from one or more...
void startDefinition()
Starts the definition of this tag declaration.
Exposes information about the current target.
TargetOptions & getTargetOpts() const
Retrieve the target options.
const llvm::Triple & getTriple() const
Returns the target triple of the primary target.
StringRef getPlatformName() const
Retrieve the name of the platform as it is used in the availability attribute.
VersionTuple getPlatformMinVersion() const
Retrieve the minimum desired version of the platform, to which the program should be compiled.
std::string HLSLEntry
The entry point name for HLSL shader being compiled as specified by -E.
A convenient class for passing around template argument information.
void addArgument(const TemplateArgumentLoc &Loc)
The base class of all kinds of template declarations (e.g., class, function, etc.).
Stores a list of template parameters for a TemplateDecl and its derived classes.
The top declaration context.
SourceLocation getBeginLoc() const
Get the begin source location.
A container of type source information.
TypeLoc getTypeLoc() const
Return the TypeLoc wrapper for the type source info.
The base class of the type hierarchy.
bool isBooleanType() const
bool isIncompleteArrayType() const
CXXRecordDecl * getAsCXXRecordDecl() const
Retrieves the CXXRecordDecl that this type refers to, either because the type is a RecordType or beca...
bool isConstantArrayType() const
bool hasIntegerRepresentation() const
Determine whether this type has an integer representation of some sort, e.g., it is an integer type o...
CXXRecordDecl * castAsCXXRecordDecl() const
bool isArithmeticType() const
bool isConstantMatrixType() const
bool isHLSLBuiltinIntangibleType() const
bool isPointerType() const
CanQualType getCanonicalTypeUnqualified() const
bool isIntegerType() const
isIntegerType() does not include complex integers (a GCC extension).
const T * castAs() const
Member-template castAs<specific type>.
bool isReferenceType() const
bool isHLSLIntangibleType() const
bool isEnumeralType() const
bool isScalarType() const
bool isIntegralType(const ASTContext &Ctx) const
Determine whether this type is an integral type.
const Type * getArrayElementTypeNoTypeQual() const
If this is an array type, return the element type of the array, potentially with type qualifiers miss...
QualType getPointeeType() const
If this is a pointer, ObjC object pointer, or block pointer, this returns the respective pointee.
bool hasUnsignedIntegerRepresentation() const
Determine whether this type has an unsigned integer representation of some sort, e....
bool isDependentType() const
Whether this type is a dependent type, meaning that its definition somehow depends on a template para...
bool isAggregateType() const
Determines whether the type is a C++ aggregate type or C aggregate or union type.
ScalarTypeKind getScalarTypeKind() const
Given that this is a scalar type, classify it.
bool hasSignedIntegerRepresentation() const
Determine whether this type has an signed integer representation of some sort, e.g....
bool isMatrixType() const
bool isHLSLResourceRecord() const
bool hasFloatingRepresentation() const
Determine whether this type has a floating-point representation of some sort, e.g....
bool isVectorType() const
bool isRealFloatingType() const
Floating point categories.
bool isHLSLAttributedResourceType() const
bool isFloatingType() const
const T * getAs() const
Member-template getAs<specific type>'.
const Type * getUnqualifiedDesugaredType() const
Return the specified type with any "sugar" removed from the type, removing any typedefs,...
bool isRecordType() const
bool isHLSLResourceRecordArray() const
void setType(QualType newType)
Represents a variable declaration or definition.
static VarDecl * Create(ASTContext &C, DeclContext *DC, SourceLocation StartLoc, SourceLocation IdLoc, const IdentifierInfo *Id, QualType T, TypeSourceInfo *TInfo, StorageClass S)
void setInitStyle(InitializationStyle Style)
@ CallInit
Call-style initialization (C++98)
void setStorageClass(StorageClass SC)
bool hasGlobalStorage() const
Returns true for all variables that do not have local storage.
StorageClass getStorageClass() const
Returns the storage class as written in the source.
Represents a GCC generic vector type.
unsigned getNumElements() const
QualType getElementType() const
void pushName(llvm::StringRef N)
void pushArrayIndex(uint64_t Index)
void pushBaseName(llvm::StringRef N)
IdentifierInfo * getNameAsIdentifier(ASTContext &AST) const
Defines the clang::TargetInfo interface.
uint32_t getResourceDimensions(llvm::dxil::ResourceDimension Dim)
bool hasCounterHandle(const CXXRecordDecl *RD)
The JSON file list parser is used to communicate input to InstallAPI.
bool isa(CodeGen::Address addr)
if(T->getSizeExpr()) TRY_TO(TraverseStmt(const_cast< Expr * >(T -> getSizeExpr())))
static bool CheckFloatOrHalfRepresentation(Sema *S, SourceLocation Loc, int ArgOrdinal, clang::QualType PassedType)
@ ICIS_NoInit
No in-class initializer.
@ TemplateName
The identifier is a template name. FIXME: Add an annotation for that.
@ OK_Ordinary
An ordinary object is located at an address in memory.
static bool CheckAllArgTypesAreCorrect(Sema *S, CallExpr *TheCall, llvm::ArrayRef< llvm::function_ref< bool(Sema *, SourceLocation, int, QualType)> > Checks)
@ AANT_ArgumentIdentifier
@ Result
The result type of a method or function.
@ Ordinary
This parameter uses ordinary ABI rules for its type.
llvm::Expected< QualType > ExpectedType
@ Template
We are parsing a template declaration.
LLVM_READONLY bool isDigit(unsigned char c)
Return true if this character is an ASCII digit: [0-9].
static bool CheckAllArgsHaveSameType(Sema *S, CallExpr *TheCall)
@ Type
The name was classified as a type.
LangAS
Defines the address space values used by the address space qualifier of QualType.
bool CreateHLSLAttributedResourceType(Sema &S, QualType Wrapped, ArrayRef< const Attr * > AttrList, QualType &ResType, HLSLAttributedResourceLocInfo *LocInfo=nullptr)
CastKind
CastKind - The kind of operation required for a conversion.
ExprValueKind
The categorization of expression values, currently following the C++11 scheme.
@ VK_PRValue
A pr-value expression (in the C++11 taxonomy) produces a temporary value.
@ VK_LValue
An l-value expression is a reference to an object with independent storage.
DynamicRecursiveASTVisitorBase< false > DynamicRecursiveASTVisitor
U cast(CodeGen::Address addr)
@ None
No keyword precedes the qualified type name.
ActionResult< Expr * > ExprResult
Visibility
Describes the different kinds of visibility that a declaration may have.
hash_code hash_value(const clang::dependencies::ModuleID &ID)
__DEVICE__ bool isnan(float __x)
__DEVICE__ _Tp abs(const std::complex< _Tp > &__c)
TypeSourceInfo * ContainedTyInfo
Describes how types, statements, expressions, and declarations should be printed.
unsigned getImplicitOrderID() const
void setCounterImplicitOrderID(unsigned Value) const
bool hasCounterImplicitOrderID() const
unsigned getSpace() const
bool hasImplicitOrderID() const
void setImplicitOrderID(unsigned Value) const
const SourceLocation & getLocation() const
const llvm::hlsl::rootsig::RootElement & getElement() const