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())
571 "struct is already HLSL buffer compatible");
585 LS->
addAttr(PackedAttr::CreateImplicit(AST));
589 if (
unsigned NumBases = StructDecl->
getNumBases()) {
590 assert(NumBases == 1 &&
"HLSL supports only one base type");
640 LS->
addAttr(PackedAttr::CreateImplicit(AST));
645 VarDecl *VD = dyn_cast<VarDecl>(D);
661 "host layout field for $Globals decl failed to be created");
678 uint32_t ImplicitBindingOrderID) {
680 HLSLResourceBindingAttr::CreateImplicit(S.
getASTContext(),
"",
"0", {});
681 Attr->setBinding(RT, std::nullopt, 0);
682 Attr->setImplicitBindingOrderID(ImplicitBindingOrderID);
689 BufDecl->setRBraceLoc(RBrace);
706 BufDecl->isCBuffer() ? RegisterType::CBuffer
716 int X,
int Y,
int Z) {
717 if (HLSLNumThreadsAttr *NT = D->
getAttr<HLSLNumThreadsAttr>()) {
718 if (NT->getX() !=
X || NT->getY() != Y || NT->getZ() != Z) {
719 Diag(NT->getLocation(), diag::err_hlsl_attribute_param_mismatch) << AL;
720 Diag(AL.
getLoc(), diag::note_conflicting_attribute);
730 int Min,
int Max,
int Preferred,
731 int SpelledArgsCount) {
732 if (HLSLWaveSizeAttr *WS = D->
getAttr<HLSLWaveSizeAttr>()) {
733 if (WS->getMin() !=
Min || WS->getMax() !=
Max ||
734 WS->getPreferred() != Preferred ||
735 WS->getSpelledArgsCount() != SpelledArgsCount) {
736 Diag(WS->getLocation(), diag::err_hlsl_attribute_param_mismatch) << AL;
737 Diag(AL.
getLoc(), diag::note_conflicting_attribute);
743 Result->setSpelledArgsCount(SpelledArgsCount);
747HLSLVkConstantIdAttr *
753 Diag(AL.
getLoc(), diag::warn_attribute_ignored) << AL;
761 Diag(VD->getLocation(), diag::err_specialization_const);
765 if (!VD->getType().isConstQualified()) {
766 Diag(VD->getLocation(), diag::err_specialization_const);
770 if (HLSLVkConstantIdAttr *CI = D->
getAttr<HLSLVkConstantIdAttr>()) {
771 if (CI->getId() != Id) {
772 Diag(CI->getLocation(), diag::err_hlsl_attribute_param_mismatch) << AL;
773 Diag(AL.
getLoc(), diag::note_conflicting_attribute);
778 HLSLVkConstantIdAttr *
Result =
785 llvm::Triple::EnvironmentType ShaderType) {
786 if (HLSLShaderAttr *NT = D->
getAttr<HLSLShaderAttr>()) {
787 if (NT->getType() != ShaderType) {
788 Diag(NT->getLocation(), diag::err_hlsl_attribute_param_mismatch) << AL;
789 Diag(AL.
getLoc(), diag::note_conflicting_attribute);
793 return HLSLShaderAttr::Create(
getASTContext(), ShaderType, AL);
796HLSLParamModifierAttr *
798 HLSLParamModifierAttr::Spelling Spelling) {
801 if (HLSLParamModifierAttr *PA = D->
getAttr<HLSLParamModifierAttr>()) {
802 if ((PA->isIn() && Spelling == HLSLParamModifierAttr::Keyword_out) ||
803 (PA->isOut() && Spelling == HLSLParamModifierAttr::Keyword_in)) {
804 D->
dropAttr<HLSLParamModifierAttr>();
806 return HLSLParamModifierAttr::Create(
808 HLSLParamModifierAttr::Keyword_inout);
810 Diag(AL.
getLoc(), diag::err_hlsl_duplicate_parameter_modifier) << AL;
811 Diag(PA->getLocation(), diag::note_conflicting_attribute);
837 if (HLSLShaderAttr::isValidShaderType(Env) && Env != llvm::Triple::Library) {
838 if (
const auto *Shader = FD->
getAttr<HLSLShaderAttr>()) {
841 if (Shader->getType() != Env) {
842 Diag(Shader->getLocation(), diag::err_hlsl_entry_shader_attr_mismatch)
854 case llvm::Triple::UnknownEnvironment:
855 case llvm::Triple::Library:
857 case llvm::Triple::RootSignature:
858 llvm_unreachable(
"rootsig environment has no functions");
860 llvm_unreachable(
"Unhandled environment in triple");
866 HLSLAppliedSemanticAttr *Semantic,
871 const auto *ShaderAttr = FD->
getAttr<HLSLShaderAttr>();
872 assert(ShaderAttr &&
"Entry point has no shader attribute");
873 llvm::Triple::EnvironmentType ST = ShaderAttr->getType();
874 auto SemanticName = Semantic->getSemanticName().upper();
879 if (SemanticName ==
"SV_POSITION") {
880 return (ST == llvm::Triple::Vertex && !IsInput) ||
881 (ST == llvm::Triple::Pixel && IsInput);
883 if (SemanticName ==
"SV_VERTEXID")
889bool SemaHLSL::determineActiveSemanticOnScalar(
FunctionDecl *FD,
892 SemanticInfo &ActiveSemantic,
893 SemaHLSL::SemanticContext &SC) {
894 if (ActiveSemantic.Semantic ==
nullptr) {
895 ActiveSemantic.Semantic = D->
getAttr<HLSLParsedSemanticAttr>();
896 if (ActiveSemantic.Semantic)
897 ActiveSemantic.Index = ActiveSemantic.Semantic->getSemanticIndex();
900 if (!ActiveSemantic.Semantic) {
906 HLSLAppliedSemanticAttr(
getASTContext(), *ActiveSemantic.Semantic,
907 ActiveSemantic.Semantic->getAttrName()->getName(),
908 ActiveSemantic.Index.value_or(0));
912 checkSemanticAnnotation(FD, D, A, SC);
913 OutputDecl->addAttr(A);
915 unsigned Location = ActiveSemantic.Index.value_or(0);
918 SC.CurrentIOType & IOType::In)) {
919 bool HasVkLocation =
false;
920 if (
auto *A = D->getAttr<HLSLVkLocationAttr>()) {
921 HasVkLocation = true;
922 Location = A->getLocation();
925 if (SC.UsesExplicitVkLocations.value_or(HasVkLocation) != HasVkLocation) {
926 Diag(D->getLocation(), diag::err_hlsl_semantic_partial_explicit_indexing);
929 SC.UsesExplicitVkLocations = HasVkLocation;
932 const ConstantArrayType *AT = dyn_cast<ConstantArrayType>(D->getType());
933 unsigned ElementCount = AT ? AT->
getZExtSize() : 1;
934 ActiveSemantic.Index = Location + ElementCount;
936 Twine BaseName = Twine(ActiveSemantic.Semantic->getAttrName()->getName());
937 for (
unsigned I = 0; I < ElementCount; ++I) {
938 Twine VariableName = BaseName.concat(Twine(Location + I));
940 auto [_, Inserted] = SC.ActiveSemantics.insert(VariableName.str());
942 Diag(D->getLocation(), diag::err_hlsl_semantic_index_overlap)
943 << VariableName.str();
954 SemanticInfo &ActiveSemantic,
955 SemaHLSL::SemanticContext &SC) {
956 if (ActiveSemantic.Semantic ==
nullptr) {
957 ActiveSemantic.Semantic = D->
getAttr<HLSLParsedSemanticAttr>();
958 if (ActiveSemantic.Semantic)
959 ActiveSemantic.Index = ActiveSemantic.Semantic->getSemanticIndex();
965 const RecordType *RT = dyn_cast<RecordType>(T);
967 return determineActiveSemanticOnScalar(FD, OutputDecl, D, ActiveSemantic,
970 const RecordDecl *RD = RT->getDecl();
971 for (FieldDecl *Field : RD->
fields()) {
972 SemanticInfo Info = ActiveSemantic;
973 if (!determineActiveSemantic(FD, OutputDecl, Field, Info, SC)) {
974 Diag(
Field->getLocation(), diag::note_hlsl_semantic_used_here) <<
Field;
977 if (ActiveSemantic.Semantic)
978 ActiveSemantic = Info;
985 const auto *ShaderAttr = FD->
getAttr<HLSLShaderAttr>();
986 assert(ShaderAttr &&
"Entry point has no shader attribute");
987 llvm::Triple::EnvironmentType ST = ShaderAttr->getType();
991 case llvm::Triple::Pixel:
992 case llvm::Triple::Vertex:
993 case llvm::Triple::Geometry:
994 case llvm::Triple::Hull:
995 case llvm::Triple::Domain:
996 case llvm::Triple::RayGeneration:
997 case llvm::Triple::Intersection:
998 case llvm::Triple::AnyHit:
999 case llvm::Triple::ClosestHit:
1000 case llvm::Triple::Miss:
1001 case llvm::Triple::Callable:
1002 if (
const auto *NT = FD->
getAttr<HLSLNumThreadsAttr>()) {
1003 diagnoseAttrStageMismatch(NT, ST,
1004 {llvm::Triple::Compute,
1005 llvm::Triple::Amplification,
1006 llvm::Triple::Mesh});
1009 if (
const auto *WS = FD->
getAttr<HLSLWaveSizeAttr>()) {
1010 diagnoseAttrStageMismatch(WS, ST,
1011 {llvm::Triple::Compute,
1012 llvm::Triple::Amplification,
1013 llvm::Triple::Mesh});
1018 case llvm::Triple::Compute:
1019 case llvm::Triple::Amplification:
1020 case llvm::Triple::Mesh:
1021 if (!FD->
hasAttr<HLSLNumThreadsAttr>()) {
1023 << llvm::Triple::getEnvironmentTypeName(ST);
1026 if (
const auto *WS = FD->
getAttr<HLSLWaveSizeAttr>()) {
1027 if (Ver < VersionTuple(6, 6)) {
1028 Diag(WS->getLocation(), diag::err_hlsl_attribute_in_wrong_shader_model)
1031 }
else if (WS->getSpelledArgsCount() > 1 && Ver < VersionTuple(6, 8)) {
1034 diag::err_hlsl_attribute_number_arguments_insufficient_shader_model)
1035 << WS << WS->getSpelledArgsCount() <<
"6.8";
1040 case llvm::Triple::RootSignature:
1041 llvm_unreachable(
"rootsig environment has no function entry point");
1043 llvm_unreachable(
"Unhandled environment in triple");
1046 SemaHLSL::SemanticContext InputSC = {};
1047 InputSC.CurrentIOType = IOType::In;
1050 SemanticInfo ActiveSemantic;
1051 ActiveSemantic.Semantic = Param->getAttr<HLSLParsedSemanticAttr>();
1052 if (ActiveSemantic.Semantic)
1053 ActiveSemantic.Index = ActiveSemantic.Semantic->getSemanticIndex();
1056 if (!determineActiveSemantic(FD, Param, Param, ActiveSemantic, InputSC)) {
1057 Diag(Param->getLocation(), diag::note_previous_decl) << Param;
1062 SemanticInfo ActiveSemantic;
1063 SemaHLSL::SemanticContext OutputSC = {};
1064 OutputSC.CurrentIOType = IOType::Out;
1065 ActiveSemantic.Semantic = FD->
getAttr<HLSLParsedSemanticAttr>();
1066 if (ActiveSemantic.Semantic)
1067 ActiveSemantic.Index = ActiveSemantic.Semantic->getSemanticIndex();
1069 determineActiveSemantic(FD, FD, FD, ActiveSemantic, OutputSC);
1072void SemaHLSL::checkSemanticAnnotation(
1074 const HLSLAppliedSemanticAttr *SemanticAttr,
const SemanticContext &SC) {
1075 auto *ShaderAttr = EntryPoint->
getAttr<HLSLShaderAttr>();
1076 assert(ShaderAttr &&
"Entry point has no shader attribute");
1077 llvm::Triple::EnvironmentType ST = ShaderAttr->getType();
1079 auto SemanticName = SemanticAttr->getSemanticName().upper();
1080 if (SemanticName ==
"SV_DISPATCHTHREADID" ||
1081 SemanticName ==
"SV_GROUPINDEX" || SemanticName ==
"SV_GROUPTHREADID" ||
1082 SemanticName ==
"SV_GROUPID") {
1084 if (ST != llvm::Triple::Compute)
1085 diagnoseSemanticStageMismatch(SemanticAttr, ST, SC.CurrentIOType,
1086 {{llvm::Triple::Compute, IOType::In}});
1088 if (SemanticAttr->getSemanticIndex() != 0) {
1089 std::string PrettyName =
1090 "'" + SemanticAttr->getSemanticName().str() +
"'";
1091 Diag(SemanticAttr->getLoc(),
1092 diag::err_hlsl_semantic_indexing_not_supported)
1098 if (SemanticName ==
"SV_POSITION") {
1101 diagnoseSemanticStageMismatch(SemanticAttr, ST, SC.CurrentIOType,
1102 {{llvm::Triple::Vertex, IOType::InOut},
1103 {llvm::Triple::Pixel, IOType::In}});
1106 if (SemanticName ==
"SV_VERTEXID") {
1107 diagnoseSemanticStageMismatch(SemanticAttr, ST, SC.CurrentIOType,
1108 {{llvm::Triple::Vertex, IOType::In}});
1112 if (SemanticName ==
"SV_TARGET") {
1113 diagnoseSemanticStageMismatch(SemanticAttr, ST, SC.CurrentIOType,
1114 {{llvm::Triple::Pixel, IOType::Out}});
1120 if (SemanticAttr->getAttrName()->getName().starts_with_insensitive(
"SV_"))
1121 llvm_unreachable(
"Unknown SemanticAttr");
1124void SemaHLSL::diagnoseAttrStageMismatch(
1125 const Attr *A, llvm::Triple::EnvironmentType Stage,
1126 std::initializer_list<llvm::Triple::EnvironmentType> AllowedStages) {
1127 SmallVector<StringRef, 8> StageStrings;
1128 llvm::transform(AllowedStages, std::back_inserter(StageStrings),
1129 [](llvm::Triple::EnvironmentType ST) {
1131 HLSLShaderAttr::ConvertEnvironmentTypeToStr(ST));
1133 Diag(A->
getLoc(), diag::err_hlsl_attr_unsupported_in_stage)
1134 << A->
getAttrName() << llvm::Triple::getEnvironmentTypeName(Stage)
1135 << (AllowedStages.size() != 1) << join(StageStrings,
", ");
1138void SemaHLSL::diagnoseSemanticStageMismatch(
1139 const Attr *A, llvm::Triple::EnvironmentType Stage, IOType CurrentIOType,
1140 std::initializer_list<SemanticStageInfo> Allowed) {
1142 for (
auto &Case : Allowed) {
1143 if (Case.Stage != Stage)
1146 if (CurrentIOType & Case.AllowedIOTypesMask)
1149 SmallVector<std::string, 8> ValidCases;
1151 Allowed, std::back_inserter(ValidCases), [](SemanticStageInfo Case) {
1152 SmallVector<std::string, 2> ValidType;
1153 if (Case.AllowedIOTypesMask & IOType::In)
1154 ValidType.push_back(
"input");
1155 if (Case.AllowedIOTypesMask & IOType::Out)
1156 ValidType.push_back(
"output");
1158 HLSLShaderAttr::ConvertEnvironmentTypeToStr(Case.Stage)) +
1159 " " + join(ValidType,
"/");
1161 Diag(A->
getLoc(), diag::err_hlsl_semantic_unsupported_iotype_for_stage)
1162 << A->
getAttrName() << (CurrentIOType & IOType::In ?
"input" :
"output")
1163 << llvm::Triple::getEnvironmentTypeName(Case.Stage)
1164 << join(ValidCases,
", ");
1168 SmallVector<StringRef, 8> StageStrings;
1170 Allowed, std::back_inserter(StageStrings), [](SemanticStageInfo Case) {
1172 HLSLShaderAttr::ConvertEnvironmentTypeToStr(Case.Stage));
1175 Diag(A->
getLoc(), diag::err_hlsl_attr_unsupported_in_stage)
1176 << A->
getAttrName() << llvm::Triple::getEnvironmentTypeName(Stage)
1177 << (Allowed.size() != 1) << join(StageStrings,
", ");
1180template <CastKind Kind>
1183 Ty = VTy->getElementType();
1188template <CastKind Kind>
1200 if (LHSFloat && RHSFloat) {
1228 if (LHSSigned == RHSSigned) {
1229 if (IsCompAssign || IntOrder >= 0)
1237 if (IntOrder != (LHSSigned ? 1 : -1)) {
1238 if (IsCompAssign || RHSSigned)
1246 if (Ctx.getIntWidth(LElTy) != Ctx.getIntWidth(RElTy)) {
1247 if (IsCompAssign || LHSSigned)
1263 QualType ElTy = Ctx.getCorrespondingUnsignedType(LHSSigned ? LElTy : RElTy);
1264 QualType NewTy = Ctx.getExtVectorType(
1274 return CK_FloatingCast;
1276 return CK_IntegralCast;
1278 return CK_IntegralToFloating;
1280 return CK_FloatingToIntegral;
1286 bool IsCompAssign) {
1293 if (!LVecTy && IsCompAssign) {
1295 RHS =
SemaRef.ImpCastExprToType(RHS.
get(), RElTy, CK_HLSLVectorTruncation);
1297 if (Ctx.hasSameUnqualifiedType(LHSType, RHSType))
1299 RHS =
SemaRef.ImpCastExprToType(RHS.
get(), LHSType,
1304 unsigned EndSz = std::numeric_limits<unsigned>::max();
1307 LSz = EndSz = LVecTy->getNumElements();
1310 assert(EndSz != std::numeric_limits<unsigned>::max() &&
1311 "one of the above should have had a value");
1315 if (IsCompAssign && LSz != EndSz) {
1317 diag::err_hlsl_vector_compound_assignment_truncation)
1318 << LHSType << RHSType;
1324 if (!IsCompAssign && LVecTy && LVecTy->getNumElements() > EndSz)
1329 if (!IsCompAssign && !LVecTy)
1333 if (Ctx.hasSameUnqualifiedType(LHSType, RHSType))
1334 return Ctx.getCommonSugaredType(LHSType, RHSType);
1342 LElTy, RElTy, IsCompAssign);
1345 "HLSL Vectors can only contain integer or floating point types");
1347 LElTy, RElTy, IsCompAssign);
1352 assert((Opc == BO_LOr || Opc == BO_LAnd) &&
1353 "Called with non-logical operator");
1355 llvm::raw_svector_ostream OS(Buff);
1357 StringRef NewFnName = Opc == BO_LOr ?
"or" :
"and";
1358 OS << NewFnName <<
"(";
1368std::pair<IdentifierInfo *, bool>
1371 std::string IdStr =
"__hlsl_rootsig_decl_" + std::to_string(Hash);
1378 return {DeclIdent,
Found};
1389 for (
auto &RootSigElement : RootElements)
1390 Elements.push_back(RootSigElement.getElement());
1394 DeclIdent,
SemaRef.getLangOpts().HLSLRootSigVer, Elements);
1396 SignatureDecl->setImplicit();
1402 if (RootSigOverrideIdent) {
1405 if (
SemaRef.LookupQualifiedName(R, DC))
1406 return dyn_cast<HLSLRootSignatureDecl>(R.getFoundDecl());
1414struct PerVisibilityBindingChecker {
1417 std::array<llvm::hlsl::BindingInfoBuilder, 8> Builders;
1421 llvm::dxbc::ShaderVisibility Vis;
1426 PerVisibilityBindingChecker(
SemaHLSL *S) : S(S) {}
1428 void trackBinding(llvm::dxbc::ShaderVisibility
Visibility,
1429 llvm::dxil::ResourceClass RC, uint32_t Space,
1430 uint32_t LowerBound, uint32_t UpperBound,
1431 const hlsl::RootSignatureElement *Elem) {
1433 assert(BuilderIndex < Builders.size() &&
1434 "Not enough builders for visibility type");
1435 Builders[BuilderIndex].trackBinding(RC, Space, LowerBound, UpperBound,
1436 static_cast<const void *
>(Elem));
1438 static_assert(llvm::to_underlying(llvm::dxbc::ShaderVisibility::All) == 0,
1439 "'All' visibility must come first");
1440 if (
Visibility == llvm::dxbc::ShaderVisibility::All)
1441 for (
size_t I = 1, E = Builders.size(); I < E; ++I)
1442 Builders[I].trackBinding(RC, Space, LowerBound, UpperBound,
1443 static_cast<const void *
>(Elem));
1445 ElemInfoMap.push_back({Elem,
Visibility,
false});
1448 ElemInfo &
getInfo(
const hlsl::RootSignatureElement *Elem) {
1449 auto It = llvm::lower_bound(
1451 [](
const auto &LHS,
const auto &RHS) {
return LHS.Elem < RHS; });
1452 assert(It->Elem == Elem &&
"Element not in map");
1456 bool checkOverlap() {
1457 llvm::sort(ElemInfoMap, [](
const auto &LHS,
const auto &RHS) {
1458 return LHS.Elem < RHS.Elem;
1461 bool HadOverlap =
false;
1463 using llvm::hlsl::BindingInfoBuilder;
1464 auto ReportOverlap = [
this,
1465 &HadOverlap](
const BindingInfoBuilder &Builder,
1466 const llvm::hlsl::Binding &Reported) {
1470 static_cast<const hlsl::RootSignatureElement *
>(Reported.Cookie);
1471 const llvm::hlsl::Binding &
Previous = Builder.findOverlapping(Reported);
1472 const auto *PrevElem =
1473 static_cast<const hlsl::RootSignatureElement *
>(
Previous.Cookie);
1475 ElemInfo &Info =
getInfo(Elem);
1480 Info.Diagnosed =
true;
1482 ElemInfo &PrevInfo =
getInfo(PrevElem);
1483 llvm::dxbc::ShaderVisibility CommonVis =
1484 Info.Vis == llvm::dxbc::ShaderVisibility::All ? PrevInfo.Vis
1487 this->S->
Diag(Elem->
getLocation(), diag::err_hlsl_resource_range_overlap)
1488 << llvm::to_underlying(Reported.RC) << Reported.LowerBound
1489 << Reported.isUnbounded() << Reported.UpperBound
1494 this->S->
Diag(PrevElem->getLocation(),
1495 diag::note_hlsl_resource_range_here);
1498 for (BindingInfoBuilder &Builder : Builders)
1499 Builder.calculateBindingInfo(ReportOverlap);
1519 bool HadError =
false;
1520 auto ReportError = [
this, &HadError](
SourceLocation Loc, uint32_t LowerBound,
1521 uint32_t UpperBound) {
1523 this->
Diag(Loc, diag::err_hlsl_invalid_rootsig_value)
1524 << LowerBound << UpperBound;
1531 this->
Diag(Loc, diag::err_hlsl_invalid_rootsig_value)
1532 << llvm::formatv(
"{0:f}", LowerBound).sstr<6>()
1533 << llvm::formatv(
"{0:f}", UpperBound).sstr<6>();
1536 auto VerifyRegister = [ReportError](
SourceLocation Loc, uint32_t Register) {
1537 if (!llvm::hlsl::rootsig::verifyRegisterValue(Register))
1538 ReportError(Loc, 0, 0xfffffffe);
1541 auto VerifySpace = [ReportError](
SourceLocation Loc, uint32_t Space) {
1542 if (!llvm::hlsl::rootsig::verifyRegisterSpace(Space))
1543 ReportError(Loc, 0, 0xffffffef);
1546 const uint32_t Version =
1547 llvm::to_underlying(
SemaRef.getLangOpts().HLSLRootSigVer);
1548 const uint32_t VersionEnum = Version - 1;
1549 auto ReportFlagError = [
this, &HadError, VersionEnum](
SourceLocation Loc) {
1551 this->
Diag(Loc, diag::err_hlsl_invalid_rootsig_flag)
1558 const llvm::hlsl::rootsig::RootElement &Elem = RootSigElem.
getElement();
1559 if (
const auto *Descriptor =
1560 std::get_if<llvm::hlsl::rootsig::RootDescriptor>(&Elem)) {
1561 VerifyRegister(Loc, Descriptor->Reg.Number);
1562 VerifySpace(Loc, Descriptor->Space);
1564 if (!llvm::hlsl::rootsig::verifyRootDescriptorFlag(Version,
1566 ReportFlagError(Loc);
1567 }
else if (
const auto *Constants =
1568 std::get_if<llvm::hlsl::rootsig::RootConstants>(&Elem)) {
1569 VerifyRegister(Loc, Constants->Reg.Number);
1570 VerifySpace(Loc, Constants->Space);
1571 }
else if (
const auto *Sampler =
1572 std::get_if<llvm::hlsl::rootsig::StaticSampler>(&Elem)) {
1573 VerifyRegister(Loc, Sampler->Reg.Number);
1574 VerifySpace(Loc, Sampler->Space);
1577 "By construction, parseFloatParam can't produce a NaN from a "
1578 "float_literal token");
1580 if (!llvm::hlsl::rootsig::verifyMaxAnisotropy(Sampler->MaxAnisotropy))
1581 ReportError(Loc, 0, 16);
1582 if (!llvm::hlsl::rootsig::verifyMipLODBias(Sampler->MipLODBias))
1583 ReportFloatError(Loc, -16.f, 15.99f);
1584 }
else if (
const auto *Clause =
1585 std::get_if<llvm::hlsl::rootsig::DescriptorTableClause>(
1587 VerifyRegister(Loc, Clause->Reg.Number);
1588 VerifySpace(Loc, Clause->Space);
1590 if (!llvm::hlsl::rootsig::verifyNumDescriptors(Clause->NumDescriptors)) {
1594 ReportError(Loc, 1, 0xfffffffe);
1597 if (!llvm::hlsl::rootsig::verifyDescriptorRangeFlag(Version, Clause->Type,
1599 ReportFlagError(Loc);
1603 PerVisibilityBindingChecker BindingChecker(
this);
1604 SmallVector<std::pair<
const llvm::hlsl::rootsig::DescriptorTableClause *,
1609 const llvm::hlsl::rootsig::RootElement &Elem = RootSigElem.
getElement();
1610 if (
const auto *Descriptor =
1611 std::get_if<llvm::hlsl::rootsig::RootDescriptor>(&Elem)) {
1612 uint32_t LowerBound(Descriptor->Reg.Number);
1613 uint32_t UpperBound(LowerBound);
1615 BindingChecker.trackBinding(
1616 Descriptor->Visibility,
1617 static_cast<llvm::dxil::ResourceClass
>(Descriptor->Type),
1618 Descriptor->Space, LowerBound, UpperBound, &RootSigElem);
1619 }
else if (
const auto *Constants =
1620 std::get_if<llvm::hlsl::rootsig::RootConstants>(&Elem)) {
1621 uint32_t LowerBound(Constants->Reg.Number);
1622 uint32_t UpperBound(LowerBound);
1624 BindingChecker.trackBinding(
1625 Constants->Visibility, llvm::dxil::ResourceClass::CBuffer,
1626 Constants->Space, LowerBound, UpperBound, &RootSigElem);
1627 }
else if (
const auto *Sampler =
1628 std::get_if<llvm::hlsl::rootsig::StaticSampler>(&Elem)) {
1629 uint32_t LowerBound(Sampler->Reg.Number);
1630 uint32_t UpperBound(LowerBound);
1632 BindingChecker.trackBinding(
1633 Sampler->Visibility, llvm::dxil::ResourceClass::Sampler,
1634 Sampler->Space, LowerBound, UpperBound, &RootSigElem);
1635 }
else if (
const auto *Clause =
1636 std::get_if<llvm::hlsl::rootsig::DescriptorTableClause>(
1639 UnboundClauses.emplace_back(Clause, &RootSigElem);
1640 }
else if (
const auto *Table =
1641 std::get_if<llvm::hlsl::rootsig::DescriptorTable>(&Elem)) {
1642 assert(UnboundClauses.size() == Table->NumClauses &&
1643 "Number of unbound elements must match the number of clauses");
1644 bool HasAnySampler =
false;
1645 bool HasAnyNonSampler =
false;
1646 uint64_t Offset = 0;
1647 bool IsPrevUnbound =
false;
1648 for (
const auto &[Clause, ClauseElem] : UnboundClauses) {
1650 if (Clause->Type == llvm::dxil::ResourceClass::Sampler)
1651 HasAnySampler =
true;
1653 HasAnyNonSampler =
true;
1655 if (HasAnySampler && HasAnyNonSampler)
1656 Diag(Loc, diag::err_hlsl_invalid_mixed_resources);
1661 if (Clause->NumDescriptors == 0)
1665 Clause->Offset == llvm::hlsl::rootsig::DescriptorTableOffsetAppend;
1667 Offset = Clause->Offset;
1669 uint64_t RangeBound = llvm::hlsl::rootsig::computeRangeBound(
1670 Offset, Clause->NumDescriptors);
1672 if (IsPrevUnbound && IsAppending)
1673 Diag(Loc, diag::err_hlsl_appending_onto_unbound);
1674 else if (!llvm::hlsl::rootsig::verifyNoOverflowedOffset(RangeBound))
1675 Diag(Loc, diag::err_hlsl_offset_overflow) << Offset << RangeBound;
1678 Offset = RangeBound + 1;
1679 IsPrevUnbound = Clause->NumDescriptors ==
1680 llvm::hlsl::rootsig::NumDescriptorsUnbounded;
1683 uint32_t LowerBound(Clause->Reg.Number);
1684 uint32_t UpperBound = llvm::hlsl::rootsig::computeRangeBound(
1685 LowerBound, Clause->NumDescriptors);
1687 BindingChecker.trackBinding(
1689 static_cast<llvm::dxil::ResourceClass
>(Clause->Type), Clause->Space,
1690 LowerBound, UpperBound, ClauseElem);
1692 UnboundClauses.clear();
1696 return BindingChecker.checkOverlap();
1701 Diag(AL.
getLoc(), diag::err_attribute_wrong_number_arguments) << AL << 1;
1706 if (
auto *RS = D->
getAttr<RootSignatureAttr>()) {
1707 if (RS->getSignatureIdent() != Ident) {
1708 Diag(AL.
getLoc(), diag::err_disallowed_duplicate_attribute) << RS;
1712 Diag(AL.
getLoc(), diag::warn_duplicate_attribute_exact) << RS;
1718 if (
auto *SignatureDecl =
1719 dyn_cast<HLSLRootSignatureDecl>(R.getFoundDecl())) {
1726 llvm::VersionTuple SMVersion =
1731 uint32_t ZMax = 1024;
1732 uint32_t ThreadMax = 1024;
1733 if (IsDXIL && SMVersion.getMajor() <= 4) {
1736 }
else if (IsDXIL && SMVersion.getMajor() == 5) {
1746 diag::err_hlsl_numthreads_argument_oor)
1755 diag::err_hlsl_numthreads_argument_oor)
1764 diag::err_hlsl_numthreads_argument_oor)
1769 if (
X * Y * Z > ThreadMax) {
1770 Diag(AL.
getLoc(), diag::err_hlsl_numthreads_invalid) << ThreadMax;
1787 if (SpelledArgsCount == 0 || SpelledArgsCount > 3)
1795 if (SpelledArgsCount > 1 &&
1799 uint32_t Preferred = 0;
1800 if (SpelledArgsCount > 2 &&
1804 if (SpelledArgsCount > 2) {
1807 diag::err_attribute_power_of_two_in_range)
1808 << AL << llvm::dxil::MinWaveSize << llvm::dxil::MaxWaveSize
1813 if (Preferred < Min || Preferred >
Max) {
1815 diag::err_attribute_power_of_two_in_range)
1816 << AL <<
Min <<
Max << Preferred;
1819 }
else if (SpelledArgsCount > 1) {
1822 diag::err_attribute_power_of_two_in_range)
1823 << AL << llvm::dxil::MinWaveSize << llvm::dxil::MaxWaveSize <<
Max;
1827 Diag(AL.
getLoc(), diag::err_attribute_argument_invalid) << AL << 1;
1830 Diag(AL.
getLoc(), diag::warn_attr_min_eq_max) << AL;
1835 diag::err_attribute_power_of_two_in_range)
1836 << AL << llvm::dxil::MinWaveSize << llvm::dxil::MaxWaveSize <<
Min;
1841 HLSLWaveSizeAttr *NewAttr =
1878 uint32_t Binding = 0;
1902 if (!T->hasUnsignedIntegerRepresentation() ||
1903 (VT && VT->getNumElements() > 3)) {
1904 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type)
1905 << AL <<
"uint/uint2/uint3";
1914 if (!T->hasFloatingRepresentation() || (VT && VT->getNumElements() > 4)) {
1915 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type)
1916 << AL <<
"float/float1/float2/float3/float4";
1924 std::optional<unsigned> Index) {
1928 QualType ValueType = VD->getType();
1929 if (
auto *FD = dyn_cast<FunctionDecl>(D))
1932 bool IsOutput =
false;
1933 if (HLSLParamModifierAttr *MA = D->
getAttr<HLSLParamModifierAttr>()) {
1940 if (SemanticName ==
"SV_DISPATCHTHREADID") {
1943 Diag(AL.
getLoc(), diag::err_hlsl_semantic_output_not_supported) << AL;
1944 if (Index.has_value())
1945 Diag(AL.
getLoc(), diag::err_hlsl_semantic_indexing_not_supported) << AL;
1950 if (SemanticName ==
"SV_GROUPINDEX") {
1952 Diag(AL.
getLoc(), diag::err_hlsl_semantic_output_not_supported) << AL;
1953 if (Index.has_value())
1954 Diag(AL.
getLoc(), diag::err_hlsl_semantic_indexing_not_supported) << AL;
1959 if (SemanticName ==
"SV_GROUPTHREADID") {
1962 Diag(AL.
getLoc(), diag::err_hlsl_semantic_output_not_supported) << AL;
1963 if (Index.has_value())
1964 Diag(AL.
getLoc(), diag::err_hlsl_semantic_indexing_not_supported) << AL;
1969 if (SemanticName ==
"SV_GROUPID") {
1972 Diag(AL.
getLoc(), diag::err_hlsl_semantic_output_not_supported) << AL;
1973 if (Index.has_value())
1974 Diag(AL.
getLoc(), diag::err_hlsl_semantic_indexing_not_supported) << AL;
1979 if (SemanticName ==
"SV_POSITION") {
1980 const auto *VT = ValueType->getAs<
VectorType>();
1981 if (!ValueType->hasFloatingRepresentation() ||
1982 (VT && VT->getNumElements() > 4))
1983 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type)
1984 << AL <<
"float/float1/float2/float3/float4";
1989 if (SemanticName ==
"SV_VERTEXID") {
1990 uint64_t SizeInBits =
SemaRef.Context.getTypeSize(ValueType);
1991 if (!ValueType->isUnsignedIntegerType() || SizeInBits != 32)
1992 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type) << AL <<
"uint";
1997 if (SemanticName ==
"SV_TARGET") {
1998 const auto *VT = ValueType->getAs<
VectorType>();
1999 if (!ValueType->hasFloatingRepresentation() ||
2000 (VT && VT->getNumElements() > 4))
2001 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type)
2002 << AL <<
"float/float1/float2/float3/float4";
2007 Diag(AL.
getLoc(), diag::err_hlsl_unknown_semantic) << AL;
2011 uint32_t IndexValue(0), ExplicitIndex(0);
2014 assert(0 &&
"HLSLUnparsedSemantic is expected to have 2 int arguments.");
2016 assert(IndexValue > 0 ? ExplicitIndex :
true);
2017 std::optional<unsigned> Index =
2018 ExplicitIndex ? std::optional<unsigned>(IndexValue) : std::nullopt;
2028 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_ast_node)
2029 << AL <<
"shader constant in a constant buffer";
2033 uint32_t SubComponent;
2043 bool IsAggregateTy = (T->isArrayType() || T->isStructureType());
2048 if (IsAggregateTy) {
2049 Diag(AL.
getLoc(), diag::err_hlsl_invalid_register_or_packoffset);
2053 if ((Component * 32 + Size) > 128) {
2054 Diag(AL.
getLoc(), diag::err_hlsl_packoffset_cross_reg_boundary);
2059 EltTy = VT->getElementType();
2061 if (Align > 32 && Component == 1) {
2064 Diag(AL.
getLoc(), diag::err_hlsl_packoffset_alignment_mismatch)
2078 if (!
SemaRef.checkStringLiteralArgumentAttr(AL, 0, Str, &ArgLoc))
2081 llvm::Triple::EnvironmentType ShaderType;
2082 if (!HLSLShaderAttr::ConvertStrToEnvironmentType(Str, ShaderType)) {
2083 Diag(AL.
getLoc(), diag::warn_attribute_type_not_supported)
2084 << AL << Str << ArgLoc;
2098 assert(AttrList.size() &&
"expected list of resource attributes");
2105 HLSLAttributedResourceType::Attributes ResAttrs;
2107 bool HasResourceClass =
false;
2108 bool HasResourceDimension =
false;
2109 for (
const Attr *A : AttrList) {
2114 case attr::HLSLResourceClass: {
2116 if (HasResourceClass) {
2118 ? diag::warn_duplicate_attribute_exact
2119 : diag::warn_duplicate_attribute)
2123 ResAttrs.ResourceClass = RC;
2124 HasResourceClass =
true;
2127 case attr::HLSLResourceDimension: {
2128 llvm::dxil::ResourceDimension RD =
2130 if (HasResourceDimension) {
2132 ? diag::warn_duplicate_attribute_exact
2133 : diag::warn_duplicate_attribute)
2137 ResAttrs.ResourceDimension = RD;
2138 HasResourceDimension =
true;
2142 if (ResAttrs.IsROV) {
2146 ResAttrs.IsROV =
true;
2148 case attr::HLSLRawBuffer:
2149 if (ResAttrs.RawBuffer) {
2153 ResAttrs.RawBuffer =
true;
2155 case attr::HLSLIsCounter:
2156 if (ResAttrs.IsCounter) {
2160 ResAttrs.IsCounter =
true;
2162 case attr::HLSLContainedType: {
2165 if (!ContainedTy.
isNull()) {
2167 ? diag::warn_duplicate_attribute_exact
2168 : diag::warn_duplicate_attribute)
2177 llvm_unreachable(
"unhandled resource attribute type");
2181 if (!HasResourceClass) {
2182 S.
Diag(AttrList.back()->getRange().getEnd(),
2183 diag::err_hlsl_missing_resource_class);
2188 Wrapped, ContainedTy, ResAttrs);
2190 if (LocInfo && ContainedTyInfo) {
2203 if (!T->isHLSLResourceType()) {
2204 Diag(AL.
getLoc(), diag::err_hlsl_attribute_needs_intangible_type)
2219 AttributeCommonInfo::AS_CXX11, 0, false ,
2224 case ParsedAttr::AT_HLSLResourceClass: {
2226 Diag(AL.
getLoc(), diag::err_attribute_argument_type)
2237 if (!HLSLResourceClassAttr::ConvertStrToResourceClass(Identifier, RC)) {
2238 Diag(ArgLoc, diag::warn_attribute_type_not_supported)
2239 <<
"ResourceClass" << Identifier;
2242 A = HLSLResourceClassAttr::Create(
getASTContext(), RC, ACI);
2246 case ParsedAttr::AT_HLSLResourceDimension: {
2247 StringRef Identifier;
2249 if (!
SemaRef.checkStringLiteralArgumentAttr(AL, 0, Identifier, &ArgLoc))
2253 llvm::dxil::ResourceDimension RD;
2254 if (!HLSLResourceDimensionAttr::ConvertStrToResourceDimension(Identifier,
2256 Diag(ArgLoc, diag::warn_attribute_type_not_supported)
2257 <<
"ResourceDimension" << Identifier;
2260 A = HLSLResourceDimensionAttr::Create(
getASTContext(), RD, ACI);
2264 case ParsedAttr::AT_HLSLROV:
2268 case ParsedAttr::AT_HLSLRawBuffer:
2272 case ParsedAttr::AT_HLSLIsCounter:
2276 case ParsedAttr::AT_HLSLContainedType: {
2278 Diag(AL.
getLoc(), diag::err_attribute_wrong_number_arguments) << AL << 1;
2284 assert(TSI &&
"no type source info for attribute argument");
2286 diag::err_incomplete_type))
2288 A = HLSLContainedTypeAttr::Create(
getASTContext(), TSI, ACI);
2293 llvm_unreachable(
"unhandled HLSL attribute");
2296 HLSLResourcesTypeAttrs.emplace_back(A);
2302 if (!HLSLResourcesTypeAttrs.size())
2308 HLSLResourcesTypeAttrs, QT, &LocInfo)) {
2309 const HLSLAttributedResourceType *RT =
2316 LocsForHLSLAttributedResources.insert(std::pair(RT, LocInfo));
2318 HLSLResourcesTypeAttrs.clear();
2326 auto I = LocsForHLSLAttributedResources.find(RT);
2327 if (I != LocsForHLSLAttributedResources.end()) {
2328 LocInfo = I->second;
2329 LocsForHLSLAttributedResources.erase(I);
2338void SemaHLSL::collectResourceBindingsOnUserRecordDecl(
const VarDecl *VD,
2339 const RecordType *RT) {
2347 "incomplete arrays inside user defined types are not supported");
2356 if (
const HLSLAttributedResourceType *AttrResType =
2357 HLSLAttributedResourceType::findHandleTypeOnResource(Ty)) {
2362 Bindings.addDeclBindingInfo(VD, RC);
2363 }
else if (
const RecordType *RT = dyn_cast<RecordType>(Ty)) {
2369 collectResourceBindingsOnUserRecordDecl(VD, RT);
2381 bool SpecifiedSpace) {
2382 int RegTypeNum =
static_cast<int>(RegType);
2385 if (D->
hasAttr<HLSLGroupSharedAddressSpaceAttr>()) {
2386 S.
Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
2391 if (
HLSLBufferDecl *CBufferOrTBuffer = dyn_cast<HLSLBufferDecl>(D)) {
2392 ResourceClass RC = CBufferOrTBuffer->isCBuffer() ? ResourceClass::CBuffer
2393 : ResourceClass::SRV;
2403 assert(
isa<VarDecl>(D) &&
"D is expected to be VarDecl or HLSLBufferDecl");
2407 if (
const HLSLAttributedResourceType *AttrResType =
2408 HLSLAttributedResourceType::findHandleTypeOnResource(
2425 if (SpecifiedSpace && !DeclaredInCOrTBuffer)
2426 S.
Diag(ArgLoc, diag::err_hlsl_space_on_global_constant);
2431 if (RegType == RegisterType::CBuffer)
2432 S.
Diag(ArgLoc, diag::warn_hlsl_deprecated_register_type_b);
2433 else if (RegType != RegisterType::C)
2434 S.
Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
2438 if (RegType == RegisterType::C)
2439 S.
Diag(ArgLoc, diag::warn_hlsl_register_type_c_packoffset);
2441 S.
Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
2451 S.
Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
2459 bool RegisterTypesDetected[5] = {
false};
2460 RegisterTypesDetected[
static_cast<int>(regType)] =
true;
2463 if (HLSLResourceBindingAttr *
attr =
2464 dyn_cast<HLSLResourceBindingAttr>(*it)) {
2467 if (RegisterTypesDetected[
static_cast<int>(otherRegType)]) {
2468 int otherRegTypeNum =
static_cast<int>(otherRegType);
2470 diag::err_hlsl_duplicate_register_annotation)
2474 RegisterTypesDetected[
static_cast<int>(otherRegType)] =
true;
2482 bool SpecifiedSpace) {
2487 "expecting VarDecl or HLSLBufferDecl");
2499 const uint64_t &Limit,
2502 uint64_t ArrayCount = 1) {
2507 if (StartSlot > Limit)
2511 if (
const auto *AT = dyn_cast<ArrayType>(T)) {
2514 if (
const auto *CAT = dyn_cast<ConstantArrayType>(AT))
2515 Count = CAT->
getSize().getZExtValue();
2519 ArrayCount * Count);
2523 if (
auto ResTy = dyn_cast<HLSLAttributedResourceType>(T)) {
2526 if (ResTy->getAttrs().ResourceClass != ResClass)
2530 uint64_t EndSlot = StartSlot + ArrayCount - 1;
2531 if (EndSlot > Limit)
2535 StartSlot = EndSlot + 1;
2540 if (
const auto *RT = dyn_cast<RecordType>(T)) {
2543 if (
const auto *CXXRD = dyn_cast<CXXRecordDecl>(RD)) {
2546 ResClass, Ctx, ArrayCount))
2553 ResClass, Ctx, ArrayCount))
2567 const uint64_t Limit = UINT32_MAX;
2568 if (SlotNum > Limit)
2573 if (RegTy == RegisterType::C || RegTy == RegisterType::I)
2576 if (
VarDecl *VD = dyn_cast<VarDecl>(TheDecl)) {
2577 uint64_t BaseSlot = SlotNum;
2585 return (BaseSlot > Limit);
2592 return (SlotNum > Limit);
2595 llvm_unreachable(
"unexpected decl type");
2599 if (
VarDecl *VD = dyn_cast<VarDecl>(TheDecl)) {
2601 if (
const auto *IAT = dyn_cast<IncompleteArrayType>(Ty))
2602 Ty = IAT->getElementType();
2604 diag::err_incomplete_type))
2608 StringRef Slot =
"";
2609 StringRef Space =
"";
2613 Diag(AL.
getLoc(), diag::err_attribute_argument_type)
2623 Diag(AL.
getLoc(), diag::err_attribute_argument_type)
2629 SpaceLoc = Loc->
getLoc();
2632 if (Str.starts_with(
"space")) {
2634 SpaceLoc = Loc->
getLoc();
2643 std::optional<unsigned> SlotNum;
2644 unsigned SpaceNum = 0;
2647 if (!Slot.empty()) {
2649 Diag(SlotLoc, diag::err_hlsl_binding_type_invalid) << Slot.substr(0, 1);
2652 if (RegType == RegisterType::I) {
2653 Diag(SlotLoc, diag::warn_hlsl_deprecated_register_type_i);
2656 const StringRef SlotNumStr = Slot.substr(1);
2661 if (SlotNumStr.getAsInteger(10, N)) {
2662 Diag(SlotLoc, diag::err_hlsl_unsupported_register_number);
2670 Diag(SlotLoc, diag::err_hlsl_register_number_too_large);
2679 if (!Space.starts_with(
"space")) {
2680 Diag(SpaceLoc, diag::err_hlsl_expected_space) << Space;
2683 StringRef SpaceNumStr = Space.substr(5);
2684 if (SpaceNumStr.getAsInteger(10, SpaceNum)) {
2685 Diag(SpaceLoc, diag::err_hlsl_expected_space) << Space;
2690 if (SlotNum.has_value())
2695 HLSLResourceBindingAttr *NewAttr =
2696 HLSLResourceBindingAttr::Create(
getASTContext(), Slot, Space, AL);
2698 NewAttr->setBinding(RegType, SlotNum, SpaceNum);
2724 while (
const auto *AT = Cur->
getAs<AttributedType>()) {
2726 if (K == attr::HLSLRowMajor || K == attr::HLSLColumnMajor) {
2730 Cur = AT->getModifiedType();
2741 ? attr::HLSLRowMajor
2742 : attr::HLSLColumnMajor;
2747 Diag(AL.
getLoc(), diag::err_hlsl_matrix_layout_non_matrix)
2756 if (ExistingKind == AttrK) {
2757 Diag(AL.
getLoc(), diag::warn_duplicate_attribute_exact)
2759 Diag(AL.
getLoc(), diag::note_previous_attribute);
2763 ExistingKind == attr::HLSLRowMajor ?
"row_major" :
"column_major");
2764 Diag(AL.
getLoc(), diag::err_hlsl_matrix_layout_conflict)
2766 Diag(AL.
getLoc(), diag::note_conflicting_attribute);
2771 if (AttrK == attr::HLSLRowMajor)
2772 return ::new (Ctx) HLSLRowMajorAttr(Ctx, AL);
2773 return ::new (Ctx) HLSLColumnMajorAttr(Ctx, AL);
2784 if (K != attr::HLSLRowMajor && K != attr::HLSLColumnMajor)
2786 if (T.isNull() || T->isDependentType())
2791 K == attr::HLSLRowMajor ?
"row_major" :
"column_major");
2792 Diag(Loc, diag::err_hlsl_matrix_layout_non_matrix) << II;
2838 llvm::DenseMap<const FunctionDecl *, unsigned> ScannedDecls;
2842 llvm::Triple::EnvironmentType CurrentShaderEnvironment;
2843 unsigned CurrentShaderStageBit;
2848 bool ReportOnlyShaderStageIssues;
2851 void SetShaderStageContext(llvm::Triple::EnvironmentType ShaderType) {
2852 static_assert(
sizeof(
unsigned) >= 4);
2853 assert(HLSLShaderAttr::isValidShaderType(ShaderType));
2854 assert((
unsigned)(ShaderType - llvm::Triple::Pixel) < 31 &&
2855 "ShaderType is too big for this bitmap");
2858 unsigned bitmapIndex = ShaderType - llvm::Triple::Pixel;
2859 CurrentShaderEnvironment = ShaderType;
2860 CurrentShaderStageBit = (1 << bitmapIndex);
2863 void SetUnknownShaderStageContext() {
2864 CurrentShaderEnvironment = llvm::Triple::UnknownEnvironment;
2865 CurrentShaderStageBit = (1 << 31);
2868 llvm::Triple::EnvironmentType GetCurrentShaderEnvironment()
const {
2869 return CurrentShaderEnvironment;
2872 bool InUnknownShaderStageContext()
const {
2873 return CurrentShaderEnvironment == llvm::Triple::UnknownEnvironment;
2877 void AddToScannedFunctions(
const FunctionDecl *FD) {
2878 unsigned &ScannedStages = ScannedDecls[FD];
2879 ScannedStages |= CurrentShaderStageBit;
2882 unsigned GetScannedStages(
const FunctionDecl *FD) {
return ScannedDecls[FD]; }
2884 bool WasAlreadyScannedInCurrentStage(
const FunctionDecl *FD) {
2885 return WasAlreadyScannedInCurrentStage(GetScannedStages(FD));
2888 bool WasAlreadyScannedInCurrentStage(
unsigned ScannerStages) {
2889 return ScannerStages & CurrentShaderStageBit;
2892 static bool NeverBeenScanned(
unsigned ScannedStages) {
2893 return ScannedStages == 0;
2897 void HandleFunctionOrMethodRef(FunctionDecl *FD, Expr *RefExpr);
2898 void CheckDeclAvailability(NamedDecl *D,
const AvailabilityAttr *AA,
2900 const AvailabilityAttr *FindAvailabilityAttr(
const Decl *D);
2901 bool HasMatchingEnvironmentOrNone(
const AvailabilityAttr *AA);
2904 DiagnoseHLSLAvailability(Sema &SemaRef)
2906 CurrentShaderEnvironment(llvm::Triple::UnknownEnvironment),
2907 CurrentShaderStageBit(0), ReportOnlyShaderStageIssues(
false) {}
2910 void RunOnTranslationUnit(
const TranslationUnitDecl *TU);
2911 void RunOnFunction(
const FunctionDecl *FD);
2913 bool VisitDeclRefExpr(DeclRefExpr *DRE)
override {
2914 FunctionDecl *FD = llvm::dyn_cast<FunctionDecl>(DRE->
getDecl());
2916 HandleFunctionOrMethodRef(FD, DRE);
2920 bool VisitMemberExpr(MemberExpr *ME)
override {
2921 FunctionDecl *FD = llvm::dyn_cast<FunctionDecl>(ME->
getMemberDecl());
2923 HandleFunctionOrMethodRef(FD, ME);
2928void DiagnoseHLSLAvailability::HandleFunctionOrMethodRef(
FunctionDecl *FD,
2931 "expected DeclRefExpr or MemberExpr");
2935 if (FD->
hasBody(FDWithBody)) {
2936 if (!WasAlreadyScannedInCurrentStage(FDWithBody))
2937 DeclsToScan.push_back(FDWithBody);
2942 const AvailabilityAttr *AA = FindAvailabilityAttr(FD);
2944 CheckDeclAvailability(
2948void DiagnoseHLSLAvailability::RunOnTranslationUnit(
2957 DeclContextsToScan.push_back(TU);
2959 while (!DeclContextsToScan.empty()) {
2960 const DeclContext *DC = DeclContextsToScan.pop_back_val();
2961 for (
auto &D : DC->
decls()) {
2968 if (llvm::dyn_cast<NamespaceDecl>(D) || llvm::dyn_cast<ExportDecl>(D)) {
2969 DeclContextsToScan.push_back(llvm::dyn_cast<DeclContext>(D));
2974 const FunctionDecl *FD = llvm::dyn_cast<FunctionDecl>(D);
2979 if (HLSLShaderAttr *ShaderAttr = FD->
getAttr<HLSLShaderAttr>()) {
2980 SetShaderStageContext(ShaderAttr->getType());
2989 for (
const auto *Redecl : FD->
redecls()) {
2990 if (Redecl->isInExportDeclContext()) {
2997 SetUnknownShaderStageContext();
3005void DiagnoseHLSLAvailability::RunOnFunction(
const FunctionDecl *FD) {
3006 assert(DeclsToScan.empty() &&
"DeclsToScan should be empty");
3007 DeclsToScan.push_back(FD);
3009 while (!DeclsToScan.empty()) {
3017 const unsigned ScannedStages = GetScannedStages(FD);
3018 if (WasAlreadyScannedInCurrentStage(ScannedStages))
3021 ReportOnlyShaderStageIssues = !NeverBeenScanned(ScannedStages);
3023 AddToScannedFunctions(FD);
3028bool DiagnoseHLSLAvailability::HasMatchingEnvironmentOrNone(
3029 const AvailabilityAttr *AA) {
3034 llvm::Triple::EnvironmentType CurrentEnv = GetCurrentShaderEnvironment();
3035 if (CurrentEnv == llvm::Triple::UnknownEnvironment)
3038 llvm::Triple::EnvironmentType AttrEnv =
3039 AvailabilityAttr::getEnvironmentType(IIEnvironment->
getName());
3041 return CurrentEnv == AttrEnv;
3044const AvailabilityAttr *
3045DiagnoseHLSLAvailability::FindAvailabilityAttr(
const Decl *D) {
3046 AvailabilityAttr
const *PartialMatch =
nullptr;
3050 for (
const auto *A : D->
attrs()) {
3051 if (
const auto *Avail = dyn_cast<AvailabilityAttr>(A)) {
3052 const AvailabilityAttr *EffectiveAvail = Avail->getEffectiveAttr();
3053 StringRef AttrPlatform = EffectiveAvail->getPlatform()->getName();
3054 StringRef TargetPlatform =
3058 if (AttrPlatform == TargetPlatform) {
3060 if (HasMatchingEnvironmentOrNone(EffectiveAvail))
3062 PartialMatch = Avail;
3066 return PartialMatch;
3071void DiagnoseHLSLAvailability::CheckDeclAvailability(
NamedDecl *D,
3072 const AvailabilityAttr *AA,
3091 if (ReportOnlyShaderStageIssues)
3097 if (InUnknownShaderStageContext())
3102 bool EnvironmentMatches = HasMatchingEnvironmentOrNone(AA);
3103 VersionTuple Introduced = AA->getIntroduced();
3112 llvm::StringRef PlatformName(
3115 llvm::StringRef CurrentEnvStr =
3116 llvm::Triple::getEnvironmentTypeName(GetCurrentShaderEnvironment());
3118 llvm::StringRef AttrEnvStr =
3119 AA->getEnvironment() ? AA->getEnvironment()->getName() :
"";
3120 bool UseEnvironment = !AttrEnvStr.empty();
3122 if (EnvironmentMatches) {
3123 SemaRef.
Diag(
Range.getBegin(), diag::warn_hlsl_availability)
3124 <<
Range << D << PlatformName << Introduced.getAsString()
3125 << UseEnvironment << CurrentEnvStr;
3127 SemaRef.
Diag(
Range.getBegin(), diag::warn_hlsl_availability_unavailable)
3131 SemaRef.
Diag(D->
getLocation(), diag::note_partial_availability_specified_here)
3132 << D << PlatformName << Introduced.getAsString()
3134 << UseEnvironment << AttrEnvStr << CurrentEnvStr;
3141 if (!DefaultCBufferDecls.empty()) {
3144 DefaultCBufferDecls);
3147 SemaRef.getCurLexicalContext()->addDecl(DefaultCBuffer);
3151 for (
const Decl *VD : DefaultCBufferDecls) {
3152 const HLSLResourceBindingAttr *RBA =
3153 VD->
getAttr<HLSLResourceBindingAttr>();
3154 if (RBA && RBA->hasRegisterSlot() &&
3155 RBA->getRegisterType() == HLSLResourceBindingAttr::RegisterType::C) {
3162 SemaRef.Consumer.HandleTopLevelDecl(DG);
3164 diagnoseAvailabilityViolations(TU);
3173 "expected member expr to have resource record type or array of them");
3179 const Expr *NonConstIndexExpr =
nullptr;
3182 if (
const DeclRefExpr *DRE = dyn_cast<DeclRefExpr>(E)) {
3183 if (!NonConstIndexExpr)
3191 diag::err_hlsl_resource_member_array_access_not_constant);
3195 if (
const auto *ASE = dyn_cast<ArraySubscriptExpr>(E)) {
3196 const Expr *IdxExpr = ASE->getIdx();
3198 NonConstIndexExpr = IdxExpr;
3200 }
else if (
const auto *SubME = dyn_cast<MemberExpr>(E)) {
3201 E = SubME->getBase();
3202 }
else if (
const auto *ICE = dyn_cast<ImplicitCastExpr>(E)) {
3203 E = ICE->getSubExpr();
3205 llvm_unreachable(
"unexpected expr type in resource member access");
3214 SemaRef.Context.getCanonicalType(
SemaRef.Context.getAddrSpaceQualType(
3217 SemaRef.Context.getLValueReferenceType(AddrSpaceType));
3220 SemaRef.Context.DeclarationNames.getCXXConversionFunctionName(
3224 [[maybe_unused]]
bool LookupSucceeded =
3225 SemaRef.LookupQualifiedName(ConvR, RD);
3226 assert(LookupSucceeded);
3235std::optional<ExprResult>
3238 const HLSLAttributedResourceType *ResTy =
3239 HLSLAttributedResourceType::findHandleTypeOnResource(
3240 BaseType.getTypePtr());
3242 ResTy->getAttrs().ResourceClass != llvm::dxil::ResourceClass::CBuffer)
3243 return std::nullopt;
3245 QualType TemplateType = ResTy->getContainedType();
3249 assert(NamedConversionDecl &&
3250 "Could not find conversion function for ConstantBuffer.");
3251 auto *ConversionDecl =
3254 return SemaRef.BuildCXXMemberCallExpr(BaseExpr.
get(), NamedConversionDecl,
3266 TI.
getTriple().getEnvironment() != llvm::Triple::EnvironmentType::Library)
3269 DiagnoseHLSLAvailability(
SemaRef).RunOnTranslationUnit(TU);
3276 for (
unsigned I = 1, N = TheCall->
getNumArgs(); I < N; ++I) {
3279 S->
Diag(TheCall->
getBeginLoc(), diag::err_vec_builtin_incompatible_vector)
3304 for (
unsigned I = 0; I < TheCall->
getNumArgs(); ++I) {
3319 if (!BaseType->isFloat32Type())
3320 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3321 << ArgOrdinal << 5 << 0
3331 BaseType = VT->getElementType();
3333 BaseType = MT->getElementType();
3335 if (!BaseType->isHalfType() && !BaseType->isFloat32Type())
3336 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3337 << ArgOrdinal << 5 << 0
3351 if (!BaseType->isDoubleType()) {
3354 return S->
Diag(Loc, diag::err_builtin_requires_double_type)
3355 << ArgOrdinal << PassedType;
3362 unsigned ArgIndex) {
3363 auto *Arg = TheCall->
getArg(ArgIndex);
3365 if (Arg->IgnoreCasts()->isModifiableLvalue(S->
Context, &OrigLoc) ==
3368 S->
Diag(OrigLoc, diag::error_hlsl_inout_lvalue) << Arg << 0;
3378 if (VecTy->getElementType()->isDoubleType())
3379 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3380 << ArgOrdinal << 1 << 0 << 1
3390 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3391 << ArgOrdinal << 5 << 1
3400 if (VecTy->getElementType()->isUnsignedIntegerType())
3403 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3404 << ArgOrdinal << 4 << 3 << 0
3413 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3414 << ArgOrdinal << 5 << 3
3420 unsigned ArgOrdinal,
unsigned Width) {
3423 ArgTy = VTy->getElementType();
3425 uint64_t ElementBitCount =
3427 if (ElementBitCount != Width) {
3429 diag::err_integer_incorrect_bit_count)
3430 << Width << ElementBitCount;
3447 unsigned ArgIndex) {
3456 diag::err_typecheck_expect_scalar_or_vector)
3457 << ArgType << Scalar;
3464 QualType Scalar,
unsigned ArgIndex) {
3475 if (
const auto *VTy = ArgType->getAs<
VectorType>()) {
3488 diag::err_typecheck_expect_scalar_or_vector_or_matrix)
3489 << ArgType << Scalar;
3494 unsigned ArgIndex) {
3499 if (!(ArgType->isScalarType() ||
3500 (VTy && VTy->getElementType()->isScalarType()))) {
3502 diag::err_typecheck_expect_any_scalar_or_vector)
3512 unsigned ArgIndex) {
3514 assert(ArgIndex < TheCall->getNumArgs());
3522 diag::err_typecheck_expect_any_scalar_or_vector)
3547 diag::err_typecheck_call_different_arg_types)
3566 Arg1ScalarTy = VTy->getElementType();
3570 Arg2ScalarTy = VTy->getElementType();
3573 S->
Diag(Arg1->
getBeginLoc(), diag::err_hlsl_builtin_scalar_vector_mismatch)
3574 << 1 << TheCall->
getCallee() << Arg1Ty << Arg2Ty;
3584 if (Arg1Length > 0 && Arg0Length != Arg1Length) {
3586 diag::err_typecheck_vector_lengths_not_equal)
3592 if (Arg2Length > 0 && Arg0Length != Arg2Length) {
3594 diag::err_typecheck_vector_lengths_not_equal)
3606 assert(TheCall->
getNumArgs() > IndexArgIndex &&
"Index argument missing");
3609 unsigned int ActualDim = 1;
3611 ActualDim = VTy->getNumElements();
3612 IndexTy = VTy->getElementType();
3616 diag::err_typecheck_expect_int)
3622 const HLSLAttributedResourceType *ResTy =
3624 assert(ResTy &&
"Resource argument must be a resource");
3625 HLSLAttributedResourceType::Attributes ResAttrs = ResTy->getAttrs();
3627 unsigned int ExpectedDim = 1;
3628 if (ResAttrs.ResourceDimension != llvm::dxil::ResourceDimension::Unknown)
3631 if (ActualDim != ExpectedDim) {
3633 diag::err_hlsl_builtin_resource_coordinate_dimension_mismatch)
3644 llvm::function_ref<
bool(
const HLSLAttributedResourceType *ResType)> Check =
3648 const HLSLAttributedResourceType *ResTy =
3652 diag::err_typecheck_expect_hlsl_resource)
3656 if (Check && Check(ResTy)) {
3658 diag::err_invalid_hlsl_resource_type)
3666 QualType BaseType,
unsigned ExpectedCount,
3668 unsigned PassedCount = 1;
3670 PassedCount = VecTy->getNumElements();
3672 if (PassedCount != ExpectedCount) {
3675 S->
Diag(Loc, diag::err_typecheck_convert_incompatible)
3687 [](
const HLSLAttributedResourceType *ResType) {
3688 return ResType->getAttrs().ResourceDimension ==
3689 llvm::dxil::ResourceDimension::Unknown;
3695 [](
const HLSLAttributedResourceType *ResType) {
3696 return ResType->getAttrs().ResourceClass !=
3697 llvm::hlsl::ResourceClass::Sampler;
3705 unsigned ExpectedDim =
3733 unsigned NextIdx = 3;
3739 diag::err_typecheck_convert_incompatible)
3747 Expr *ComponentArg = TheCall->
getArg(NextIdx);
3751 diag::err_typecheck_convert_incompatible)
3758 std::optional<llvm::APSInt> ComponentOpt =
3761 int64_t ComponentVal = ComponentOpt->getSExtValue();
3762 if (ComponentVal != 0) {
3765 assert(ComponentVal >= 0 && ComponentVal <= 3 &&
3766 "The component is not in the expected range.");
3768 diag::err_hlsl_gathercmp_invalid_component)
3778 const HLSLAttributedResourceType *ResourceTy =
3781 unsigned ExpectedDim =
3790 assert(ResourceTy->hasContainedType() &&
3791 "Expecting a contained type for resource with a dimension "
3793 QualType ReturnType = ResourceTy->getContainedType();
3797 S.
Diag(TheCall->
getBeginLoc(), diag::err_hlsl_samplecmp_requires_float);
3803 ReturnType = VecTy->getElementType();
3816 [](
const HLSLAttributedResourceType *ResType) {
3817 return ResType->getAttrs().ResourceDimension ==
3818 llvm::dxil::ResourceDimension::Unknown;
3826 unsigned ExpectedDim =
3835 EltTy = VTy->getElementType();
3850 TheCall->
setType(ResourceTy->getContainedType());
3855 unsigned MinArgs, MaxArgs;
3883 const HLSLAttributedResourceType *ResourceTy =
3885 unsigned ExpectedDim =
3888 unsigned NextIdx = 3;
3897 diag::err_typecheck_convert_incompatible)
3932 diag::err_typecheck_convert_incompatible)
3938 assert(ResourceTy->hasContainedType() &&
3939 "Expecting a contained type for resource with a dimension "
3941 QualType ReturnType = ResourceTy->getContainedType();
3944 S.
Diag(TheCall->
getBeginLoc(), diag::err_hlsl_samplecmp_requires_float);
3957 switch (BuiltinID) {
3958 case Builtin::BI__builtin_hlsl_adduint64: {
3959 if (
SemaRef.checkArgCount(TheCall, 2))
3973 if (NumElementsArg != 2 && NumElementsArg != 4) {
3975 << 1 << 64 << NumElementsArg * 32;
3989 case Builtin::BI__builtin_hlsl_resource_getpointer: {
3990 if (
SemaRef.checkArgCountRange(TheCall, 1, 2) ||
3997 QualType ContainedTy = ResourceTy->getContainedType();
3998 auto ReturnType =
SemaRef.Context.getAddrSpaceQualType(
4001 ReturnType =
SemaRef.Context.getPointerType(ReturnType);
4006 case Builtin::BI__builtin_hlsl_resource_getpointer_typed: {
4007 if (
SemaRef.checkArgCount(TheCall, 3) ||
4014 "expected pointer type for second argument");
4021 diag::err_invalid_use_of_array_type);
4025 auto ReturnType =
SemaRef.Context.getAddrSpaceQualType(
4028 ReturnType =
SemaRef.Context.getPointerType(ReturnType);
4033 case Builtin::BI__builtin_hlsl_resource_load_with_status: {
4034 if (
SemaRef.checkArgCount(TheCall, 3) ||
4037 SemaRef.getASTContext().UnsignedIntTy) ||
4039 SemaRef.getASTContext().UnsignedIntTy) ||
4045 QualType ReturnType = ResourceTy->getContainedType();
4050 case Builtin::BI__builtin_hlsl_resource_load_with_status_typed: {
4051 if (
SemaRef.checkArgCount(TheCall, 4) ||
4054 SemaRef.getASTContext().UnsignedIntTy) ||
4056 SemaRef.getASTContext().UnsignedIntTy) ||
4062 "expected pointer type for second argument");
4069 diag::err_invalid_use_of_array_type);
4075 case Builtin::BI__builtin_hlsl_resource_load_level:
4077 case Builtin::BI__builtin_hlsl_resource_sample:
4079 case Builtin::BI__builtin_hlsl_resource_sample_bias:
4081 case Builtin::BI__builtin_hlsl_resource_sample_grad:
4083 case Builtin::BI__builtin_hlsl_resource_sample_level:
4085 case Builtin::BI__builtin_hlsl_resource_sample_cmp:
4087 case Builtin::BI__builtin_hlsl_resource_sample_cmp_level_zero:
4089 case Builtin::BI__builtin_hlsl_resource_calculate_lod:
4090 case Builtin::BI__builtin_hlsl_resource_calculate_lod_unclamped:
4092 case Builtin::BI__builtin_hlsl_resource_gather:
4094 case Builtin::BI__builtin_hlsl_resource_gather_cmp:
4096 case Builtin::BI__builtin_hlsl_resource_uninitializedhandle: {
4097 assert(TheCall->
getNumArgs() == 1 &&
"expected 1 arg");
4103 case Builtin::BI__builtin_hlsl_resource_handlefrombinding: {
4104 assert(TheCall->
getNumArgs() == 6 &&
"expected 6 args");
4110 case Builtin::BI__builtin_hlsl_resource_handlefromimplicitbinding: {
4111 assert(TheCall->
getNumArgs() == 6 &&
"expected 6 args");
4117 case Builtin::BI__builtin_hlsl_resource_counterhandlefromimplicitbinding: {
4118 assert(TheCall->
getNumArgs() == 3 &&
"expected 3 args");
4121 auto *MainResType = MainHandleTy->
getAs<HLSLAttributedResourceType>();
4122 auto MainAttrs = MainResType->getAttrs();
4123 assert(!MainAttrs.IsCounter &&
"cannot create a counter from a counter");
4124 MainAttrs.IsCounter =
true;
4126 MainResType->getWrappedType(), MainResType->getContainedType(),
4130 TheCall->
setType(CounterHandleTy);
4133 case Builtin::BI__builtin_hlsl_and:
4134 case Builtin::BI__builtin_hlsl_or: {
4135 if (
SemaRef.checkArgCount(TheCall, 2))
4149 case Builtin::BI__builtin_hlsl_all:
4150 case Builtin::BI__builtin_hlsl_any: {
4151 if (
SemaRef.checkArgCount(TheCall, 1))
4157 case Builtin::BI__builtin_hlsl_asdouble: {
4158 if (
SemaRef.checkArgCount(TheCall, 2))
4162 SemaRef.Context.UnsignedIntTy,
4167 SemaRef.Context.UnsignedIntTy,
4176 case Builtin::BI__builtin_hlsl_elementwise_clamp: {
4177 if (
SemaRef.BuiltinElementwiseTernaryMath(
4183 case Builtin::BI__builtin_hlsl_dot: {
4185 if (
SemaRef.BuiltinVectorToScalarMath(TheCall))
4191 case Builtin::BI__builtin_hlsl_elementwise_firstbithigh:
4192 case Builtin::BI__builtin_hlsl_elementwise_firstbitlow: {
4193 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
4203 EltTy = VecTy->getElementType();
4204 ResTy =
SemaRef.Context.getExtVectorType(ResTy, VecTy->getNumElements());
4217 case Builtin::BI__builtin_hlsl_select: {
4218 if (
SemaRef.checkArgCount(TheCall, 3))
4226 if (VTy && VTy->getElementType()->isBooleanType() &&
4231 case Builtin::BI__builtin_hlsl_elementwise_saturate:
4232 case Builtin::BI__builtin_hlsl_elementwise_rcp: {
4233 if (
SemaRef.checkArgCount(TheCall, 1))
4239 diag::err_builtin_invalid_arg_type)
4242 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
4246 case Builtin::BI__builtin_hlsl_elementwise_degrees:
4247 case Builtin::BI__builtin_hlsl_elementwise_radians:
4248 case Builtin::BI__builtin_hlsl_elementwise_rsqrt:
4249 case Builtin::BI__builtin_hlsl_elementwise_frac:
4250 case Builtin::BI__builtin_hlsl_elementwise_ddx_coarse:
4251 case Builtin::BI__builtin_hlsl_elementwise_ddy_coarse:
4252 case Builtin::BI__builtin_hlsl_elementwise_ddx_fine:
4253 case Builtin::BI__builtin_hlsl_elementwise_ddy_fine: {
4254 if (
SemaRef.checkArgCount(TheCall, 1))
4259 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
4263 case Builtin::BI__builtin_hlsl_elementwise_isinf:
4264 case Builtin::BI__builtin_hlsl_elementwise_isnan: {
4265 if (
SemaRef.checkArgCount(TheCall, 1))
4270 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
4275 case Builtin::BI__builtin_hlsl_lerp: {
4276 if (
SemaRef.checkArgCount(TheCall, 3))
4283 if (
SemaRef.BuiltinElementwiseTernaryMath(TheCall))
4287 case Builtin::BI__builtin_hlsl_mad: {
4288 if (
SemaRef.BuiltinElementwiseTernaryMath(
4294 case Builtin::BI__builtin_hlsl_mul: {
4295 if (
SemaRef.checkArgCount(TheCall, 2))
4304 if (
const auto *VTy = T->getAs<
VectorType>())
4305 return VTy->getElementType();
4307 return MTy->getElementType();
4311 QualType EltTy0 = getElemType(Ty0);
4320 if (IsVec0 && IsMat1) {
4323 }
else if (IsMat0 && IsVec1) {
4327 assert(IsMat0 && IsMat1);
4337 case Builtin::BI__builtin_hlsl_normalize: {
4338 if (
SemaRef.checkArgCount(TheCall, 1))
4349 case Builtin::BI__builtin_elementwise_fma: {
4350 if (
SemaRef.checkArgCount(TheCall, 3) ||
4365 case Builtin::BI__builtin_hlsl_transpose: {
4366 if (
SemaRef.checkArgCount(TheCall, 1))
4375 << 1 << 3 << 0 << 0 << ArgTy;
4380 MatTy->getElementType(), MatTy->getNumColumns(), MatTy->getNumRows());
4384 case Builtin::BI__builtin_hlsl_elementwise_sign: {
4385 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
4393 case Builtin::BI__builtin_hlsl_step: {
4394 if (
SemaRef.checkArgCount(TheCall, 2))
4406 case Builtin::BI__builtin_hlsl_wave_active_all_equal: {
4407 if (
SemaRef.checkArgCount(TheCall, 1))
4421 unsigned NumElts = VecTy->getNumElements();
4431 case Builtin::BI__builtin_hlsl_wave_active_max:
4432 case Builtin::BI__builtin_hlsl_wave_active_min:
4433 case Builtin::BI__builtin_hlsl_wave_active_sum:
4434 case Builtin::BI__builtin_hlsl_wave_active_product: {
4435 if (
SemaRef.checkArgCount(TheCall, 1))
4448 case Builtin::BI__builtin_hlsl_wave_active_bit_or:
4449 case Builtin::BI__builtin_hlsl_wave_active_bit_xor:
4450 case Builtin::BI__builtin_hlsl_wave_active_bit_and: {
4451 if (
SemaRef.checkArgCount(TheCall, 1))
4466 (VTy && VTy->getElementType()->isIntegerType()))) {
4468 diag::err_builtin_invalid_arg_type)
4469 << ArgTyExpr <<
SemaRef.Context.UnsignedIntTy << 1 << 0 << 0;
4479 case Builtin::BI__builtin_elementwise_bitreverse: {
4487 case Builtin::BI__builtin_hlsl_wave_prefix_count_bits: {
4488 if (
SemaRef.checkArgCount(TheCall, 1))
4493 if (!(
ArgType->isScalarType())) {
4495 diag::err_typecheck_expect_any_scalar_or_vector)
4500 if (!(
ArgType->isBooleanType())) {
4502 diag::err_typecheck_expect_any_scalar_or_vector)
4509 case Builtin::BI__builtin_hlsl_wave_read_lane_at: {
4510 if (
SemaRef.checkArgCount(TheCall, 2))
4518 diag::err_typecheck_convert_incompatible)
4519 << ArgTyIndex <<
SemaRef.Context.UnsignedIntTy << 1 << 0 << 0;
4532 case Builtin::BI__builtin_hlsl_wave_get_lane_index: {
4533 if (
SemaRef.checkArgCount(TheCall, 0))
4537 case Builtin::BI__builtin_hlsl_wave_prefix_sum:
4538 case Builtin::BI__builtin_hlsl_wave_prefix_product: {
4539 if (
SemaRef.checkArgCount(TheCall, 1))
4552 case Builtin::BI__builtin_hlsl_quad_read_across_x:
4553 case Builtin::BI__builtin_hlsl_quad_read_across_y: {
4554 if (
SemaRef.checkArgCount(TheCall, 1))
4566 case Builtin::BI__builtin_hlsl_elementwise_splitdouble: {
4567 if (
SemaRef.checkArgCount(TheCall, 3))
4573 SemaRef.Context.UnsignedIntTy, 1) ||
4575 SemaRef.Context.UnsignedIntTy, 2))
4583 case Builtin::BI__builtin_hlsl_elementwise_clip: {
4584 if (
SemaRef.checkArgCount(TheCall, 1))
4591 case Builtin::BI__builtin_elementwise_acos:
4592 case Builtin::BI__builtin_elementwise_asin:
4593 case Builtin::BI__builtin_elementwise_atan:
4594 case Builtin::BI__builtin_elementwise_atan2:
4595 case Builtin::BI__builtin_elementwise_ceil:
4596 case Builtin::BI__builtin_elementwise_cos:
4597 case Builtin::BI__builtin_elementwise_cosh:
4598 case Builtin::BI__builtin_elementwise_exp:
4599 case Builtin::BI__builtin_elementwise_exp2:
4600 case Builtin::BI__builtin_elementwise_exp10:
4601 case Builtin::BI__builtin_elementwise_floor:
4602 case Builtin::BI__builtin_elementwise_fmod:
4603 case Builtin::BI__builtin_elementwise_log:
4604 case Builtin::BI__builtin_elementwise_log2:
4605 case Builtin::BI__builtin_elementwise_log10:
4606 case Builtin::BI__builtin_elementwise_pow:
4607 case Builtin::BI__builtin_elementwise_roundeven:
4608 case Builtin::BI__builtin_elementwise_sin:
4609 case Builtin::BI__builtin_elementwise_sinh:
4610 case Builtin::BI__builtin_elementwise_sqrt:
4611 case Builtin::BI__builtin_elementwise_tan:
4612 case Builtin::BI__builtin_elementwise_tanh:
4613 case Builtin::BI__builtin_elementwise_trunc: {
4619 case Builtin::BI__builtin_hlsl_buffer_update_counter: {
4620 assert(TheCall->
getNumArgs() == 2 &&
"expected 2 args");
4621 auto checkResTy = [](
const HLSLAttributedResourceType *ResTy) ->
bool {
4622 return !(ResTy->getAttrs().ResourceClass == ResourceClass::UAV &&
4623 ResTy->getAttrs().RawBuffer && ResTy->hasContainedType());
4628 std::optional<llvm::APSInt> Offset =
4630 if (!Offset.has_value() ||
std::abs(Offset->getExtValue()) != 1) {
4632 diag::err_hlsl_expect_arg_const_int_one_or_neg_one)
4638 case Builtin::BI__builtin_hlsl_elementwise_f16tof32: {
4639 if (
SemaRef.checkArgCount(TheCall, 1))
4650 ArgTy = VTy->getElementType();
4653 diag::err_builtin_invalid_arg_type)
4662 case Builtin::BI__builtin_hlsl_elementwise_f32tof16: {
4663 if (
SemaRef.checkArgCount(TheCall, 1))
4678 WorkList.push_back(BaseTy);
4679 while (!WorkList.empty()) {
4680 QualType T = WorkList.pop_back_val();
4681 T = T.getCanonicalType().getUnqualifiedType();
4682 if (
const auto *AT = dyn_cast<ConstantArrayType>(T)) {
4690 for (uint64_t Ct = 0; Ct < AT->
getZExtSize(); ++Ct)
4691 llvm::append_range(List, ElementFields);
4696 if (
const auto *VT = dyn_cast<VectorType>(T)) {
4697 List.insert(List.end(), VT->getNumElements(), VT->getElementType());
4700 if (
const auto *MT = dyn_cast<ConstantMatrixType>(T)) {
4701 List.insert(List.end(), MT->getNumElementsFlattened(),
4702 MT->getElementType());
4705 if (
const auto *RD = T->getAsCXXRecordDecl()) {
4706 if (RD->isStandardLayout())
4707 RD = RD->getStandardLayoutBaseWithFields();
4711 if (RD->
isUnion() || !RD->isAggregate()) {
4717 for (
const auto *FD : RD->
fields())
4718 if (!FD->isUnnamedBitField())
4719 FieldTypes.push_back(FD->
getType());
4721 std::reverse(FieldTypes.begin(), FieldTypes.end());
4722 llvm::append_range(WorkList, FieldTypes);
4726 if (!RD->isStandardLayout()) {
4728 for (
const auto &
Base : RD->bases())
4729 FieldTypes.push_back(
Base.getType());
4730 std::reverse(FieldTypes.begin(), FieldTypes.end());
4731 llvm::append_range(WorkList, FieldTypes);
4766 if (
SemaRef.Context.getTypeSize(QT) / 8 > 16)
4772 int ArraySize = VT->getNumElements();
4777 QualType ElTy = VT->getElementType();
4781 if (
SemaRef.Context.getTypeSize(QT) / 8 > 16)
4797 if (
SemaRef.getASTContext().hasSameType(T1, T2))
4806 return llvm::equal(T1Types, T2Types,
4808 return SemaRef.IsLayoutCompatible(LHS, RHS);
4817 bool HadError =
false;
4819 for (
unsigned i = 0, e =
New->getNumParams(); i != e; ++i) {
4827 const auto *NDAttr = NewParam->
getAttr<HLSLParamModifierAttr>();
4828 unsigned NSpellingIdx = (NDAttr ? NDAttr->getSpellingListIndex() : 0);
4829 const auto *ODAttr = OldParam->
getAttr<HLSLParamModifierAttr>();
4830 unsigned OSpellingIdx = (ODAttr ? ODAttr->getSpellingListIndex() : 0);
4832 if (NSpellingIdx != OSpellingIdx) {
4834 diag::err_hlsl_param_qualifier_mismatch)
4835 << NDAttr << NewParam;
4851 if (
SemaRef.getASTContext().hasSameUnqualifiedType(SrcTy, DestTy))
4866 llvm_unreachable(
"HLSL doesn't support pointers.");
4869 llvm_unreachable(
"HLSL doesn't support complex types.");
4871 llvm_unreachable(
"HLSL doesn't support fixed point types.");
4873 llvm_unreachable(
"Should have returned before this");
4883 llvm_unreachable(
"HLSL doesn't support complex types.");
4885 llvm_unreachable(
"HLSL doesn't support fixed point types.");
4890 llvm_unreachable(
"HLSL doesn't support pointers.");
4892 llvm_unreachable(
"Should have returned before this");
4898 llvm_unreachable(
"HLSL doesn't support pointers.");
4901 llvm_unreachable(
"HLSL doesn't support fixed point types.");
4905 llvm_unreachable(
"HLSL doesn't support complex types.");
4908 llvm_unreachable(
"Unhandled scalar cast");
4929 !(SrcMatTy && SrcMatTy->getNumElementsFlattened() == 1))
4935 SrcTy = SrcMatTy->getElementType();
4940 for (
unsigned I = 0, Size = DestTypes.size(); I < Size; ++I) {
4941 if (DestTypes[I]->isUnionType())
4973 if (SrcTypes.size() < DestTypes.size())
4976 unsigned SrcSize = SrcTypes.size();
4977 unsigned DstSize = DestTypes.size();
4979 for (I = 0; I < DstSize && I < SrcSize; I++) {
4980 if (SrcTypes[I]->isUnionType() || DestTypes[I]->isUnionType())
4988 for (; I < SrcSize; I++) {
4989 if (SrcTypes[I]->isUnionType())
4996 assert(Param->hasAttr<HLSLParamModifierAttr>() &&
4997 "We should not get here without a parameter modifier expression");
4998 const auto *
Attr = Param->getAttr<HLSLParamModifierAttr>();
5005 << Arg << (IsInOut ? 1 : 0);
5011 QualType Ty = Param->getType().getNonLValueExprType(Ctx);
5018 << Arg << (IsInOut ? 1 : 0);
5030 SemaRef.PerformCopyInitialization(Entity, Param->getBeginLoc(), ArgOpV);
5036 auto *OpV =
new (Ctx)
5041 Res =
SemaRef.ActOnBinOp(
SemaRef.getCurScope(), Param->getBeginLoc(),
5042 tok::equal, ArgOpV, OpV);
5058 "Pointer and reference types cannot be inout or out parameters");
5059 Ty =
SemaRef.getASTContext().getLValueReferenceType(Ty);
5075 for (
const auto *FD : RD->
fields()) {
5079 assert(RD->getNumBases() <= 1 &&
5080 "HLSL doesn't support multiple inheritance");
5081 return RD->getNumBases()
5086 if (
const auto *AT = dyn_cast<ArrayType>(Ty)) {
5087 if (
const auto *CAT = dyn_cast<ConstantArrayType>(AT))
5099 bool IsVKPushConstant = IsVulkan && VD->
hasAttr<HLSLVkPushConstantAttr>();
5104 !VD->
hasAttr<HLSLVkConstantIdAttr>() && !IsVKPushConstant &&
5110 if (
Decl->getType().hasAddressSpace())
5113 if (
Decl->getType()->isDependentType())
5125 if (
Decl->
hasAttr<HLSLVkExtBuiltinOutputAttr>()) {
5139 llvm::Triple::Vulkan;
5140 if (IsVulkan &&
Decl->
hasAttr<HLSLVkPushConstantAttr>()) {
5141 if (HasDeclaredAPushConstant)
5147 HasDeclaredAPushConstant =
true;
5174class StructBindingContext {
5177 HLSLResourceBindingAttr *RegBindingsAttrs[4];
5178 unsigned RegBindingOffset[4];
5181 static_assert(
static_cast<unsigned>(RegisterType::SRV) == 0 &&
5182 static_cast<unsigned>(RegisterType::UAV) == 1 &&
5183 static_cast<unsigned>(RegisterType::CBuffer) == 2 &&
5184 static_cast<unsigned>(RegisterType::Sampler) == 3,
5185 "unexpected register type values");
5188 HLSLVkBindingAttr *VkBindingAttr;
5189 unsigned VkBindingOffset;
5194 StructBindingContext(
VarDecl *VD) {
5195 for (
unsigned i = 0; i < 4; ++i) {
5196 RegBindingsAttrs[i] =
nullptr;
5197 RegBindingOffset[i] = 0;
5199 VkBindingAttr =
nullptr;
5200 VkBindingOffset = 0;
5206 if (
auto *RBA = dyn_cast<HLSLResourceBindingAttr>(A)) {
5208 unsigned RegTypeIdx =
static_cast<unsigned>(RegType);
5211 RegBindingsAttrs[RegTypeIdx] = RBA;
5216 if (
auto *VBA = dyn_cast<HLSLVkBindingAttr>(A))
5217 VkBindingAttr = VBA;
5224 Attr *createBindingAttr(SemaHLSL &S, ASTContext &AST,
RegisterType RegType,
5225 unsigned Range,
bool HasCounter) {
5226 assert(
static_cast<unsigned>(RegType) < 4 &&
"unexpected register type");
5228 if (VkBindingAttr) {
5229 unsigned Offset = VkBindingOffset;
5230 VkBindingOffset +=
Range;
5231 return HLSLVkBindingAttr::CreateImplicit(
5232 AST, VkBindingAttr->getBinding() + Offset, VkBindingAttr->getSet(),
5233 VkBindingAttr->getRange());
5236 HLSLResourceBindingAttr *RBA =
5237 RegBindingsAttrs[
static_cast<unsigned>(RegType)];
5238 HLSLResourceBindingAttr *NewAttr =
nullptr;
5240 if (RBA && RBA->hasRegisterSlot()) {
5243 unsigned Offset = RegBindingOffset[
static_cast<unsigned>(RegType)];
5244 RegBindingOffset[
static_cast<unsigned>(RegType)] += Range;
5246 unsigned NewSlotNumber = RBA->getSlotNumber() + Offset;
5247 StringRef NewSlotNumberStr =
5249 NewAttr = HLSLResourceBindingAttr::CreateImplicit(
5250 AST, NewSlotNumberStr, RBA->getSpace(), RBA->getRange());
5251 NewAttr->setBinding(RegType, NewSlotNumber, RBA->getSpaceNumber());
5255 NewAttr = HLSLResourceBindingAttr::CreateImplicit(AST,
"",
"0", {});
5256 NewAttr->setBinding(RegType, std::nullopt,
5257 RBA ? RBA->getSpaceNumber() : 0);
5261 NewAttr->setImplicitCounterBindingOrderID(
5270static void createGlobalResourceDeclForStruct(
5272 QualType ResTy, StructBindingContext &BindingCtx) {
5274 "expected resource type or array of resources");
5285 while (
const auto *AT = dyn_cast<ArrayType>(SingleResTy)) {
5286 const auto *CAT = dyn_cast<ConstantArrayType>(AT);
5291 const HLSLAttributedResourceType *ResHandleTy =
5292 HLSLAttributedResourceType::findHandleTypeOnResource(SingleResTy);
5296 Attr *BindingAttr = BindingCtx.createBindingAttr(
5298 ResDecl->
addAttr(BindingAttr);
5299 ResDecl->
addAttr(InternalLinkageAttr::CreateImplicit(AST));
5308 HLSLAssociatedResourceDeclAttr::CreateImplicit(AST, ResDecl));
5315static void handleArrayOfStructWithResources(
5317 EmbeddedResourceNameBuilder &NameBuilder, StructBindingContext &BindingCtx);
5322static void handleStructWithResources(
Sema &S,
VarDecl *ParentVD,
5324 EmbeddedResourceNameBuilder &NameBuilder,
5325 StructBindingContext &BindingCtx) {
5328 assert(RD->
getNumBases() <= 1 &&
"HLSL doesn't support multiple inheritance");
5335 handleStructWithResources(S, ParentVD, BaseRD, NameBuilder, BindingCtx);
5349 createGlobalResourceDeclForStruct(S, ParentVD, FD->
getLocation(), II,
5352 handleStructWithResources(S, ParentVD, RD, NameBuilder, BindingCtx);
5354 }
else if (
const auto *ArrayTy = dyn_cast<ConstantArrayType>(FDTy)) {
5356 "resource arrays should have been already handled");
5357 handleArrayOfStructWithResources(S, ParentVD, ArrayTy, NameBuilder,
5366handleArrayOfStructWithResources(
Sema &S,
VarDecl *ParentVD,
5368 EmbeddedResourceNameBuilder &NameBuilder,
5369 StructBindingContext &BindingCtx) {
5377 if (!SubCAT && !ElementRD)
5380 for (
unsigned I = 0, E = CAT->
getSize().getZExtValue(); I < E; ++I) {
5383 handleStructWithResources(S, ParentVD, ElementRD, NameBuilder,
5386 handleArrayOfStructWithResources(S, ParentVD, SubCAT, NameBuilder,
5399void SemaHLSL::handleGlobalStructOrArrayOfWithResources(
VarDecl *VD) {
5400 EmbeddedResourceNameBuilder NameBuilder(VD->
getName());
5401 StructBindingContext BindingCtx(VD);
5405 "Expected non-resource struct or array type");
5408 handleStructWithResources(
SemaRef, VD, RD, NameBuilder, BindingCtx);
5412 if (
const auto *CAT = dyn_cast<ConstantArrayType>(VDTy)) {
5413 handleArrayOfStructWithResources(
SemaRef, VD, CAT, NameBuilder, BindingCtx);
5421 if (
SemaRef.RequireCompleteType(
5424 diag::err_typecheck_decl_incomplete_type)) {
5438 DefaultCBufferDecls.push_back(VD);
5443 collectResourceBindingsOnVarDecl(VD);
5445 if (VD->
hasAttr<HLSLVkConstantIdAttr>())
5457 processExplicitBindingsOnDecl(VD);
5495 handleGlobalStructOrArrayOfWithResources(VD);
5499 if (VD->
hasAttr<HLSLGroupSharedAddressSpaceAttr>())
5508 "expected resource record type");
5524 const char *CreateMethodName;
5526 CreateMethodName = HasCounter ?
"__createFromBindingWithImplicitCounter"
5527 :
"__createFromBinding";
5529 CreateMethodName = HasCounter
5530 ?
"__createFromImplicitBindingWithImplicitCounter"
5531 :
"__createFromImplicitBinding";
5536 if (!CreateMethod) {
5541 "create method lookup should always succeed for built-in resource "
5550 Args.push_back(RegSlot);
5558 Args.push_back(OrderId);
5564 Args.push_back(Space);
5568 Args.push_back(RangeSize);
5572 Args.push_back(Index);
5574 StringRef VarName = VD->
getName();
5582 Args.push_back(NameCast);
5590 Args.push_back(CounterId);
5613 SemaRef.CheckCompleteVariableDeclaration(VD);
5619 "expected array of resource records");
5640 lookupMethod(
SemaRef, ResourceDecl,
5641 HasCounter ?
"__createFromBindingWithImplicitCounter"
5642 :
"__createFromBinding",
5646 CreateMethod = lookupMethod(
5648 HasCounter ?
"__createFromImplicitBindingWithImplicitCounter"
5649 :
"__createFromImplicitBinding",
5681std::optional<const DeclBindingInfo *> SemaHLSL::inferGlobalBinding(
Expr *E) {
5682 if (
auto *Ternary = dyn_cast<ConditionalOperator>(E)) {
5683 auto TrueInfo = inferGlobalBinding(Ternary->getTrueExpr());
5684 auto FalseInfo = inferGlobalBinding(Ternary->getFalseExpr());
5685 if (!TrueInfo || !FalseInfo)
5686 return std::nullopt;
5687 if (*TrueInfo != *FalseInfo)
5688 return std::nullopt;
5692 if (
auto *ASE = dyn_cast<ArraySubscriptExpr>(E))
5701 if (
const auto *AttrResType =
5702 HLSLAttributedResourceType::findHandleTypeOnResource(Ty)) {
5704 return Bindings.getDeclBindingInfo(VD, RC);
5711void SemaHLSL::trackLocalResource(
VarDecl *VD,
Expr *E) {
5712 std::optional<const DeclBindingInfo *> ExprBinding = inferGlobalBinding(E);
5715 diag::warn_hlsl_assigning_local_resource_is_not_unique)
5720 if (*ExprBinding ==
nullptr)
5723 auto PrevBinding = Assigns.find(VD);
5724 if (PrevBinding == Assigns.end()) {
5726 Assigns.insert({VD, *ExprBinding});
5731 if (*ExprBinding != PrevBinding->second) {
5733 diag::warn_hlsl_assigning_local_resource_is_not_unique)
5735 SemaRef.Diag(VD->getLocation(), diag::note_var_declared_here) << VD;
5746 "expected LHS to be a resource record or array of resource records");
5747 if (Opc != BO_Assign)
5752 while (
auto *ASE = dyn_cast<ArraySubscriptExpr>(E))
5760 SemaRef.Diag(Loc, diag::err_hlsl_assign_to_global_resource) << VD;
5765 trackLocalResource(VD, RHSExpr);
5773void SemaHLSL::collectResourceBindingsOnVarDecl(
VarDecl *VD) {
5775 "expected global variable that contains HLSL resource");
5778 if (
const HLSLBufferDecl *CBufferOrTBuffer = dyn_cast<HLSLBufferDecl>(VD)) {
5779 Bindings.addDeclBindingInfo(VD, CBufferOrTBuffer->isCBuffer()
5780 ? ResourceClass::CBuffer
5781 : ResourceClass::SRV);
5794 if (
const HLSLAttributedResourceType *AttrResType =
5795 HLSLAttributedResourceType::findHandleTypeOnResource(Ty)) {
5796 Bindings.addDeclBindingInfo(VD, AttrResType->getAttrs().ResourceClass);
5801 if (
const RecordType *RT = dyn_cast<RecordType>(Ty))
5802 collectResourceBindingsOnUserRecordDecl(VD, RT);
5808void SemaHLSL::processExplicitBindingsOnDecl(
VarDecl *VD) {
5811 bool HasBinding =
false;
5812 for (Attr *A : VD->
attrs()) {
5815 if (
auto PA = VD->
getAttr<HLSLVkPushConstantAttr>())
5816 Diag(PA->getLoc(), diag::err_hlsl_attr_incompatible) << A << PA;
5819 HLSLResourceBindingAttr *RBA = dyn_cast<HLSLResourceBindingAttr>(A);
5820 if (!RBA || !RBA->hasRegisterSlot())
5825 assert(RT != RegisterType::I &&
"invalid or obsolete register type should "
5826 "never have an attribute created");
5828 if (RT == RegisterType::C) {
5829 if (Bindings.hasBindingInfoForDecl(VD))
5831 diag::warn_hlsl_user_defined_type_missing_member)
5832 <<
static_cast<int>(RT);
5840 if (DeclBindingInfo *BI = Bindings.getDeclBindingInfo(VD, RC)) {
5845 diag::warn_hlsl_user_defined_type_missing_member)
5846 <<
static_cast<int>(RT);
5854class InitListTransformer {
5858 QualType *DstIt =
nullptr;
5859 Expr **ArgIt =
nullptr;
5865 bool castInitializer(Expr *E) {
5866 assert(DstIt &&
"This should always be something!");
5867 if (DstIt == DestTypes.end()) {
5869 ArgExprs.push_back(E);
5874 DstIt = DestTypes.begin();
5877 Ctx, *DstIt,
false);
5882 ArgExprs.push_back(
Init);
5887 bool buildInitializerListImpl(Expr *E) {
5889 if (
auto *
Init = dyn_cast<InitListExpr>(E)) {
5890 for (
auto *SubInit :
Init->inits())
5891 if (!buildInitializerListImpl(SubInit))
5901 return castInitializer(E);
5915 if (
auto *VecTy = Ty->
getAs<VectorType>()) {
5920 for (uint64_t I = 0; I <
Size; ++I) {
5922 SizeTy, SourceLocation());
5928 if (!castInitializer(ElExpr.
get()))
5933 if (
auto *MTy = Ty->
getAs<ConstantMatrixType>()) {
5934 unsigned Rows = MTy->getNumRows();
5935 unsigned Cols = MTy->getNumColumns();
5936 QualType ElemTy = MTy->getElementType();
5938 for (
unsigned R = 0;
R < Rows; ++
R) {
5939 for (
unsigned C = 0;
C < Cols; ++
C) {
5952 if (!castInitializer(ElExpr.
get()))
5960 if (
auto *ArrTy = dyn_cast<ConstantArrayType>(Ty.
getTypePtr())) {
5964 for (uint64_t I = 0; I <
Size; ++I) {
5966 SizeTy, SourceLocation());
5971 if (!buildInitializerListImpl(ElExpr.
get()))
5978 llvm::SmallVector<CXXRecordDecl *> RecordDecls;
5979 RecordDecls.push_back(RD);
5980 while (RecordDecls.back()->getNumBases()) {
5981 CXXRecordDecl *D = RecordDecls.back();
5983 "HLSL doesn't support multiple inheritance");
5984 RecordDecls.push_back(
5987 while (!RecordDecls.empty()) {
5988 CXXRecordDecl *RD = RecordDecls.pop_back_val();
5989 for (
auto *FD : RD->
fields()) {
5990 if (FD->isUnnamedBitField())
5998 if (!buildInitializerListImpl(Res.
get()))
6006 Expr *generateInitListsImpl(QualType Ty) {
6008 assert(ArgIt != ArgExprs.end() &&
"Something is off in iteration!");
6013 llvm::SmallVector<Expr *>
Inits;
6018 if (
auto *ATy = Ty->
getAs<VectorType>()) {
6019 ElTy = ATy->getElementType();
6020 Size = ATy->getNumElements();
6021 }
else if (
auto *CMTy = Ty->
getAs<ConstantMatrixType>()) {
6022 ElTy = CMTy->getElementType();
6023 Size = CMTy->getNumElementsFlattened();
6026 ElTy = VTy->getElementType();
6027 Size = VTy->getZExtSize();
6029 for (uint64_t I = 0; I <
Size; ++I)
6030 Inits.push_back(generateInitListsImpl(ElTy));
6033 llvm::SmallVector<CXXRecordDecl *> RecordDecls;
6034 RecordDecls.push_back(RD);
6035 while (RecordDecls.back()->getNumBases()) {
6036 CXXRecordDecl *D = RecordDecls.back();
6038 "HLSL doesn't support multiple inheritance");
6039 RecordDecls.push_back(
6042 while (!RecordDecls.empty()) {
6043 CXXRecordDecl *RD = RecordDecls.pop_back_val();
6044 for (
auto *FD : RD->
fields())
6045 if (!FD->isUnnamedBitField())
6050 new (Ctx) InitListExpr(Ctx,
Inits.front()->getBeginLoc(),
Inits,
6051 Inits.back()->getEndLoc(),
false);
6052 NewInit->setType(Ty);
6057 llvm::SmallVector<QualType, 16> DestTypes;
6058 llvm::SmallVector<Expr *, 16> ArgExprs;
6059 InitListTransformer(Sema &SemaRef,
const InitializedEntity &Entity)
6060 : S(SemaRef), Ctx(SemaRef.getASTContext()),
6061 Wrap(Entity.
getType()->isIncompleteArrayType()) {
6062 InitTy = Entity.
getType().getNonReferenceType();
6072 DstIt = DestTypes.begin();
6075 bool buildInitializerList(Expr *E) {
return buildInitializerListImpl(E); }
6077 Expr *generateInitLists() {
6078 assert(!ArgExprs.empty() &&
6079 "Call buildInitializerList to generate argument expressions.");
6080 ArgIt = ArgExprs.begin();
6082 return generateInitListsImpl(InitTy);
6083 llvm::SmallVector<Expr *>
Inits;
6084 while (ArgIt != ArgExprs.end())
6085 Inits.push_back(generateInitListsImpl(InitTy));
6088 new (Ctx) InitListExpr(Ctx,
Inits.front()->getBeginLoc(),
Inits,
6089 Inits.back()->getEndLoc(),
false);
6090 llvm::APInt ArySize(64,
Inits.size());
6092 ArraySizeModifier::Normal, 0));
6104 if (
const ArrayType *AT = dyn_cast<ArrayType>(Ty)) {
6111 if (
const auto *RT = Ty->
getAs<RecordType>()) {
6115 if (
const auto *CXXRD = dyn_cast<CXXRecordDecl>(RD)) {
6135 if (
Init->getType()->isScalarType())
6138 InitListTransformer ILT(
SemaRef, Entity);
6140 for (
unsigned I = 0; I <
Init->getNumInits(); ++I) {
6148 Init->setInit(I, E);
6150 if (!ILT.buildInitializerList(E))
6153 size_t ExpectedSize = ILT.DestTypes.size();
6154 size_t ActualSize = ILT.ArgExprs.size();
6155 if (ExpectedSize == 0 && ActualSize == 0)
6162 InitTy =
SemaRef.getASTContext().removeAddrSpaceQualType(InitTy);
6164 SemaRef.Diag(
Init->getBeginLoc(), diag::err_hlsl_incorrect_num_initializers)
6165 << (int)(ExpectedSize < ActualSize) << InitTy
6166 << ExpectedSize << ActualSize;
6176 assert(ExpectedSize > 0 &&
6177 "The expected size of an incomplete array type must be at least 1.");
6179 ((ActualSize + ExpectedSize - 1) / ExpectedSize) * ExpectedSize;
6187 InitTy =
SemaRef.getASTContext().removeAddrSpaceQualType(InitTy);
6188 if (ExpectedSize != ActualSize) {
6189 int TooManyOrFew = ActualSize > ExpectedSize ? 1 : 0;
6190 SemaRef.Diag(
Init->getBeginLoc(), diag::err_hlsl_incorrect_num_initializers)
6191 << TooManyOrFew << InitTy << ExpectedSize << ActualSize;
6198 Init->resizeInits(Ctx, NewInit->getNumInits());
6199 for (
unsigned I = 0; I < NewInit->getNumInits(); ++I)
6200 Init->updateInit(Ctx, I, NewInit->getInit(I));
6208 S.
Diag(OpLoc, diag::err_builtin_matrix_invalid_member)
6218 StringRef AccessorName = CompName->
getName();
6219 assert(!AccessorName.empty() &&
"Matrix Accessor must have a name");
6221 unsigned Rows = MT->getNumRows();
6222 unsigned Cols = MT->getNumColumns();
6223 bool IsZeroBasedAccessor =
false;
6224 unsigned ChunkLen = 0;
6225 if (AccessorName.size() < 2)
6227 "length 4 for zero based: \'_mRC\' or "
6228 "length 3 for one-based: \'_RC\' accessor",
6231 if (AccessorName[0] ==
'_') {
6232 if (AccessorName[1] ==
'm') {
6233 IsZeroBasedAccessor =
true;
6240 S, AccessorName,
"zero based: \'_mRC\' or one-based: \'_RC\' accessor",
6243 if (AccessorName.size() % ChunkLen != 0) {
6244 const llvm::StringRef
Expected = IsZeroBasedAccessor
6245 ?
"zero based: '_mRC' accessor"
6246 :
"one-based: '_RC' accessor";
6251 auto isDigit = [](
char c) {
return c >=
'0' &&
c <=
'9'; };
6252 auto isZeroBasedIndex = [](
unsigned i) {
return i <= 3; };
6253 auto isOneBasedIndex = [](
unsigned i) {
return i >= 1 && i <= 4; };
6255 bool HasRepeated =
false;
6257 unsigned NumComponents = 0;
6258 const char *Begin = AccessorName.data();
6260 for (
unsigned I = 0, E = AccessorName.size(); I < E; I += ChunkLen) {
6261 const char *Chunk = Begin + I;
6262 char RowChar = 0, ColChar = 0;
6263 if (IsZeroBasedAccessor) {
6265 if (Chunk[0] !=
'_' || Chunk[1] !=
'm') {
6266 char Bad = (Chunk[0] !=
'_') ? Chunk[0] : Chunk[1];
6268 S, StringRef(&Bad, 1),
"\'_m\' prefix",
6275 if (Chunk[0] !=
'_')
6277 S, StringRef(&Chunk[0], 1),
"\'_\' prefix",
6284 bool IsDigitsError =
false;
6286 unsigned BadPos = IsZeroBasedAccessor ? 2 : 1;
6290 IsDigitsError =
true;
6294 unsigned BadPos = IsZeroBasedAccessor ? 3 : 2;
6298 IsDigitsError =
true;
6303 unsigned Row = RowChar -
'0';
6304 unsigned Col = ColChar -
'0';
6306 bool HasIndexingError =
false;
6307 if (IsZeroBasedAccessor) {
6309 if (!isZeroBasedIndex(Row)) {
6310 S.
Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
6312 HasIndexingError =
true;
6314 if (!isZeroBasedIndex(Col)) {
6315 S.
Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
6317 HasIndexingError =
true;
6321 if (!isOneBasedIndex(Row)) {
6322 S.
Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
6324 HasIndexingError =
true;
6326 if (!isOneBasedIndex(Col)) {
6327 S.
Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
6329 HasIndexingError =
true;
6336 if (HasIndexingError)
6342 bool HasBoundsError =
false;
6344 Diag(OpLoc, diag::err_hlsl_matrix_index_out_of_bounds)
6346 HasBoundsError =
true;
6349 Diag(OpLoc, diag::err_hlsl_matrix_index_out_of_bounds)
6351 HasBoundsError =
true;
6356 unsigned FlatIndex = Row * Cols + Col;
6357 if (Seen[FlatIndex])
6359 Seen[FlatIndex] =
true;
6362 if (NumComponents == 0 || NumComponents > 4) {
6363 S.
Diag(OpLoc, diag::err_hlsl_matrix_swizzle_invalid_length)
6368 QualType ElemTy = MT->getElementType();
6369 if (NumComponents == 1)
6375 for (Sema::ExtVectorDeclsType::iterator
6379 if ((*I)->getUnderlyingType() == VT)
6390 trackLocalResource(VDecl,
Init);
6392 const HLSLVkConstantIdAttr *ConstIdAttr =
6393 VDecl->
getAttr<HLSLVkConstantIdAttr>();
6400 if (!
Init->isCXX11ConstantExpr(Context, &InitValue)) {
6410 int ConstantID = ConstIdAttr->getId();
6411 llvm::APInt IDVal(Context.getIntWidth(Context.IntTy), ConstantID);
6413 ConstIdAttr->getLocation());
6417 if (
C->getType()->getCanonicalTypeUnqualified() !=
6421 Context.getTrivialTypeSourceInfo(
6422 Init->getType(),
Init->getExprLoc()),
6441 if (!Params || Params->
size() != 1)
6454 if (
auto *TTP = dyn_cast<TemplateTypeParmDecl>(P)) {
6455 if (TTP->hasDefaultArgument()) {
6456 TemplateArgs.
addArgument(TTP->getDefaultArgument());
6459 }
else if (
auto *NTTP = dyn_cast<NonTypeTemplateParmDecl>(P)) {
6460 if (NTTP->hasDefaultArgument()) {
6461 TemplateArgs.
addArgument(NTTP->getDefaultArgument());
6464 }
else if (
auto *TTPD = dyn_cast<TemplateTemplateParmDecl>(P)) {
6465 if (TTPD->hasDefaultArgument()) {
6466 TemplateArgs.
addArgument(TTPD->getDefaultArgument());
6473 return SemaRef.CheckTemplateIdType(
6475 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 bool findExistingMatrixLayoutMarker(QualType T, attr::Kind &ExistingKind)
Walks the existing AttributedType sugar of T looking for a previously applied HLSLRowMajor/HLSLColumn...
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 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 const Type * getHostLayoutFieldType(QualType QT)
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.
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
void setInvalid(bool b=true) 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
RecordDecl * getDefinitionOrSelf() 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)
Attr * buildMatrixLayoutTypeAttr(QualType T, 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)
bool IsConstantBufferElementCompatible(QualType T1)
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)
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 handleVkConstantIdAttr(Decl *D, const ParsedAttr &AL)
Decl * ActOnStartBuffer(Scope *BufferScope, bool CBuffer, SourceLocation KwLoc, IdentifierInfo *Ident, SourceLocation IdentLoc, SourceLocation LBrace)
bool diagnoseMatrixLayoutInstantiation(attr::Kind K, QualType T, SourceLocation Loc)
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