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");
83 assert(RT !=
nullptr);
87 *RT = RegisterType::SRV;
91 *RT = RegisterType::UAV;
95 *RT = RegisterType::CBuffer;
99 *RT = RegisterType::Sampler;
103 *RT = RegisterType::C;
107 *RT = RegisterType::I;
116 case RegisterType::SRV:
118 case RegisterType::UAV:
120 case RegisterType::CBuffer:
122 case RegisterType::Sampler:
124 case RegisterType::C:
126 case RegisterType::I:
129 llvm_unreachable(
"unexpected RegisterType value");
134 case RegisterType::SRV:
135 return ResourceClass::SRV;
136 case RegisterType::UAV:
137 return ResourceClass::UAV;
138 case RegisterType::CBuffer:
139 return ResourceClass::CBuffer;
140 case RegisterType::Sampler:
141 return ResourceClass::Sampler;
142 case RegisterType::C:
143 case RegisterType::I:
147 llvm_unreachable(
"unexpected RegisterType value");
151 const auto *BT = dyn_cast<BuiltinType>(
Type);
155 return Builtin::BI__builtin_get_spirv_spec_constant_int;
158 switch (BT->getKind()) {
159 case BuiltinType::Bool:
160 return Builtin::BI__builtin_get_spirv_spec_constant_bool;
161 case BuiltinType::Short:
162 return Builtin::BI__builtin_get_spirv_spec_constant_short;
163 case BuiltinType::Int:
164 return Builtin::BI__builtin_get_spirv_spec_constant_int;
165 case BuiltinType::LongLong:
166 return Builtin::BI__builtin_get_spirv_spec_constant_longlong;
167 case BuiltinType::UShort:
168 return Builtin::BI__builtin_get_spirv_spec_constant_ushort;
169 case BuiltinType::UInt:
170 return Builtin::BI__builtin_get_spirv_spec_constant_uint;
171 case BuiltinType::ULongLong:
172 return Builtin::BI__builtin_get_spirv_spec_constant_ulonglong;
173 case BuiltinType::Half:
174 return Builtin::BI__builtin_get_spirv_spec_constant_half;
175 case BuiltinType::Float:
176 return Builtin::BI__builtin_get_spirv_spec_constant_float;
177 case BuiltinType::Double:
178 return Builtin::BI__builtin_get_spirv_spec_constant_double;
187 llvm::raw_svector_ostream OS(Buffer);
194 ResourceClass ResClass) {
196 "DeclBindingInfo already added");
202 DeclToBindingListIndex.try_emplace(VD, BindingsList.size());
203 return &BindingsList.emplace_back(VD, ResClass);
207 ResourceClass ResClass) {
208 auto Entry = DeclToBindingListIndex.find(VD);
209 if (Entry != DeclToBindingListIndex.end()) {
210 for (
unsigned Index = Entry->getSecond();
211 Index < BindingsList.size() && BindingsList[Index].Decl == VD;
213 if (BindingsList[Index].ResClass == ResClass)
214 return &BindingsList[Index];
221 return DeclToBindingListIndex.contains(VD);
233 getASTContext(), LexicalParent, CBuffer, KwLoc, Ident, IdentLoc, LBrace);
236 auto RC = CBuffer ? llvm::hlsl::ResourceClass::CBuffer
237 : llvm::hlsl::ResourceClass::SRV;
249 if (T->isArrayType() || T->isStructureType() || T->isConstantMatrixType())
256 assert(Context.getTypeSize(T) <= 64 &&
257 "Scalar bit widths larger than 64 not supported");
260 return Context.getTypeSize(T) / 8;
267 constexpr unsigned CBufferAlign = 16;
268 if (
const auto *RD = T->getAsRecordDecl()) {
270 for (
const FieldDecl *Field : RD->fields()) {
277 unsigned AlignSize = llvm::alignTo(Size, FieldAlign);
278 if ((AlignSize % CBufferAlign) + FieldSize > CBufferAlign) {
279 FieldAlign = CBufferAlign;
282 Size = llvm::alignTo(Size, FieldAlign);
289 unsigned ElementCount = AT->getSize().getZExtValue();
290 if (ElementCount == 0)
293 unsigned ElementSize =
295 unsigned AlignedElementSize = llvm::alignTo(ElementSize, CBufferAlign);
296 return AlignedElementSize * (ElementCount - 1) + ElementSize;
300 unsigned ElementCount = VT->getNumElements();
301 unsigned ElementSize =
303 return ElementSize * ElementCount;
306 return Context.getTypeSize(T) / 8;
317 bool HasPackOffset =
false;
318 bool HasNonPackOffset =
false;
320 VarDecl *Var = dyn_cast<VarDecl>(Field);
323 if (Field->hasAttr<HLSLPackOffsetAttr>()) {
324 PackOffsetVec.emplace_back(Var, Field->
getAttr<HLSLPackOffsetAttr>());
325 HasPackOffset =
true;
327 HasNonPackOffset =
true;
334 if (HasNonPackOffset)
341 std::sort(PackOffsetVec.begin(), PackOffsetVec.end(),
342 [](
const std::pair<VarDecl *, HLSLPackOffsetAttr *> &LHS,
343 const std::pair<VarDecl *, HLSLPackOffsetAttr *> &RHS) {
344 return LHS.second->getOffsetInBytes() <
345 RHS.second->getOffsetInBytes();
347 for (
unsigned i = 0; i < PackOffsetVec.size() - 1; i++) {
348 VarDecl *Var = PackOffsetVec[i].first;
349 HLSLPackOffsetAttr *
Attr = PackOffsetVec[i].second;
351 unsigned Begin =
Attr->getOffsetInBytes();
352 unsigned End = Begin + Size;
353 unsigned NextBegin = PackOffsetVec[i + 1].second->getOffsetInBytes();
354 if (End > NextBegin) {
355 VarDecl *NextVar = PackOffsetVec[i + 1].first;
367 CAT = dyn_cast<ConstantArrayType>(
369 return CAT !=
nullptr;
380static const HLSLAttributedResourceType *
383 "expected array of resource records");
385 while (
const ArrayType *AT = dyn_cast<ArrayType>(Ty))
387 return HLSLAttributedResourceType::findHandleTypeOnResource(Ty);
390static const HLSLAttributedResourceType *
404 return RD->isEmpty();
433 Base.getType()->castAsCXXRecordDecl()))
444 assert(RD ==
nullptr &&
445 "there should be at most 1 record by a given name in a scope");
462 Name.append(NameBaseII->
getName());
469 size_t NameLength = Name.size();
478 Name.append(llvm::Twine(suffix).str());
479 II = &AST.
Idents.
get(Name, tok::TokenKind::identifier);
486 Name.truncate(NameLength);
501 if (
const auto *CAT = dyn_cast<ConstantArrayType>(Ty)) {
503 S, CAT->getElementType()->getUnqualifiedDesugaredType());
508 CAT->getSizeModifier(),
509 CAT->getIndexTypeCVRQualifiers())
547 "struct is already HLSL buffer compatible");
561 LS->
addAttr(PackedAttr::CreateImplicit(AST));
565 if (
unsigned NumBases = StructDecl->
getNumBases()) {
566 assert(NumBases == 1 &&
"HLSL supports only one base type");
616 LS->
addAttr(PackedAttr::CreateImplicit(AST));
621 VarDecl *VD = dyn_cast<VarDecl>(D);
637 "host layout field for $Globals decl failed to be created");
654 uint32_t ImplicitBindingOrderID) {
656 HLSLResourceBindingAttr::CreateImplicit(S.
getASTContext(),
"",
"0", {});
657 Attr->setBinding(RT, std::nullopt, 0);
658 Attr->setImplicitBindingOrderID(ImplicitBindingOrderID);
665 BufDecl->setRBraceLoc(RBrace);
682 BufDecl->isCBuffer() ? RegisterType::CBuffer
692 int X,
int Y,
int Z) {
693 if (HLSLNumThreadsAttr *NT = D->
getAttr<HLSLNumThreadsAttr>()) {
694 if (NT->getX() !=
X || NT->getY() != Y || NT->getZ() != Z) {
695 Diag(NT->getLocation(), diag::err_hlsl_attribute_param_mismatch) << AL;
696 Diag(AL.
getLoc(), diag::note_conflicting_attribute);
706 int Min,
int Max,
int Preferred,
707 int SpelledArgsCount) {
708 if (HLSLWaveSizeAttr *WS = D->
getAttr<HLSLWaveSizeAttr>()) {
709 if (WS->getMin() !=
Min || WS->getMax() !=
Max ||
710 WS->getPreferred() != Preferred ||
711 WS->getSpelledArgsCount() != SpelledArgsCount) {
712 Diag(WS->getLocation(), diag::err_hlsl_attribute_param_mismatch) << AL;
713 Diag(AL.
getLoc(), diag::note_conflicting_attribute);
719 Result->setSpelledArgsCount(SpelledArgsCount);
723HLSLVkConstantIdAttr *
729 Diag(AL.
getLoc(), diag::warn_attribute_ignored) << AL;
737 Diag(VD->getLocation(), diag::err_specialization_const);
741 if (!VD->getType().isConstQualified()) {
742 Diag(VD->getLocation(), diag::err_specialization_const);
746 if (HLSLVkConstantIdAttr *CI = D->
getAttr<HLSLVkConstantIdAttr>()) {
747 if (CI->getId() != Id) {
748 Diag(CI->getLocation(), diag::err_hlsl_attribute_param_mismatch) << AL;
749 Diag(AL.
getLoc(), diag::note_conflicting_attribute);
754 HLSLVkConstantIdAttr *
Result =
761 llvm::Triple::EnvironmentType ShaderType) {
762 if (HLSLShaderAttr *NT = D->
getAttr<HLSLShaderAttr>()) {
763 if (NT->getType() != ShaderType) {
764 Diag(NT->getLocation(), diag::err_hlsl_attribute_param_mismatch) << AL;
765 Diag(AL.
getLoc(), diag::note_conflicting_attribute);
769 return HLSLShaderAttr::Create(
getASTContext(), ShaderType, AL);
772HLSLParamModifierAttr *
774 HLSLParamModifierAttr::Spelling Spelling) {
777 if (HLSLParamModifierAttr *PA = D->
getAttr<HLSLParamModifierAttr>()) {
778 if ((PA->isIn() && Spelling == HLSLParamModifierAttr::Keyword_out) ||
779 (PA->isOut() && Spelling == HLSLParamModifierAttr::Keyword_in)) {
780 D->
dropAttr<HLSLParamModifierAttr>();
782 return HLSLParamModifierAttr::Create(
784 HLSLParamModifierAttr::Keyword_inout);
786 Diag(AL.
getLoc(), diag::err_hlsl_duplicate_parameter_modifier) << AL;
787 Diag(PA->getLocation(), diag::note_conflicting_attribute);
813 if (HLSLShaderAttr::isValidShaderType(Env) && Env != llvm::Triple::Library) {
814 if (
const auto *Shader = FD->
getAttr<HLSLShaderAttr>()) {
817 if (Shader->getType() != Env) {
818 Diag(Shader->getLocation(), diag::err_hlsl_entry_shader_attr_mismatch)
830 case llvm::Triple::UnknownEnvironment:
831 case llvm::Triple::Library:
833 case llvm::Triple::RootSignature:
834 llvm_unreachable(
"rootsig environment has no functions");
836 llvm_unreachable(
"Unhandled environment in triple");
842 HLSLAppliedSemanticAttr *Semantic,
847 const auto *ShaderAttr = FD->
getAttr<HLSLShaderAttr>();
848 assert(ShaderAttr &&
"Entry point has no shader attribute");
849 llvm::Triple::EnvironmentType ST = ShaderAttr->getType();
850 auto SemanticName = Semantic->getSemanticName().upper();
855 if (SemanticName ==
"SV_POSITION") {
856 return (ST == llvm::Triple::Vertex && !IsInput) ||
857 (ST == llvm::Triple::Pixel && IsInput);
859 if (SemanticName ==
"SV_VERTEXID")
865bool SemaHLSL::determineActiveSemanticOnScalar(
FunctionDecl *FD,
868 SemanticInfo &ActiveSemantic,
869 SemaHLSL::SemanticContext &SC) {
870 if (ActiveSemantic.Semantic ==
nullptr) {
871 ActiveSemantic.Semantic = D->
getAttr<HLSLParsedSemanticAttr>();
872 if (ActiveSemantic.Semantic)
873 ActiveSemantic.Index = ActiveSemantic.Semantic->getSemanticIndex();
876 if (!ActiveSemantic.Semantic) {
882 HLSLAppliedSemanticAttr(
getASTContext(), *ActiveSemantic.Semantic,
883 ActiveSemantic.Semantic->getAttrName()->getName(),
884 ActiveSemantic.Index.value_or(0));
888 checkSemanticAnnotation(FD, D, A, SC);
889 OutputDecl->addAttr(A);
891 unsigned Location = ActiveSemantic.Index.value_or(0);
894 SC.CurrentIOType & IOType::In)) {
895 bool HasVkLocation =
false;
896 if (
auto *A = D->getAttr<HLSLVkLocationAttr>()) {
897 HasVkLocation = true;
898 Location = A->getLocation();
901 if (SC.UsesExplicitVkLocations.value_or(HasVkLocation) != HasVkLocation) {
902 Diag(D->getLocation(), diag::err_hlsl_semantic_partial_explicit_indexing);
905 SC.UsesExplicitVkLocations = HasVkLocation;
908 const ConstantArrayType *AT = dyn_cast<ConstantArrayType>(D->getType());
909 unsigned ElementCount = AT ? AT->
getZExtSize() : 1;
910 ActiveSemantic.Index = Location + ElementCount;
912 Twine BaseName = Twine(ActiveSemantic.Semantic->getAttrName()->getName());
913 for (
unsigned I = 0; I < ElementCount; ++I) {
914 Twine VariableName = BaseName.concat(Twine(Location + I));
916 auto [_, Inserted] = SC.ActiveSemantics.insert(VariableName.str());
918 Diag(D->getLocation(), diag::err_hlsl_semantic_index_overlap)
919 << VariableName.str();
930 SemanticInfo &ActiveSemantic,
931 SemaHLSL::SemanticContext &SC) {
932 if (ActiveSemantic.Semantic ==
nullptr) {
933 ActiveSemantic.Semantic = D->
getAttr<HLSLParsedSemanticAttr>();
934 if (ActiveSemantic.Semantic)
935 ActiveSemantic.Index = ActiveSemantic.Semantic->getSemanticIndex();
941 const RecordType *RT = dyn_cast<RecordType>(T);
943 return determineActiveSemanticOnScalar(FD, OutputDecl, D, ActiveSemantic,
946 const RecordDecl *RD = RT->getDecl();
947 for (FieldDecl *Field : RD->
fields()) {
948 SemanticInfo Info = ActiveSemantic;
949 if (!determineActiveSemantic(FD, OutputDecl, Field, Info, SC)) {
950 Diag(
Field->getLocation(), diag::note_hlsl_semantic_used_here) <<
Field;
953 if (ActiveSemantic.Semantic)
954 ActiveSemantic = Info;
961 const auto *ShaderAttr = FD->
getAttr<HLSLShaderAttr>();
962 assert(ShaderAttr &&
"Entry point has no shader attribute");
963 llvm::Triple::EnvironmentType ST = ShaderAttr->getType();
967 case llvm::Triple::Pixel:
968 case llvm::Triple::Vertex:
969 case llvm::Triple::Geometry:
970 case llvm::Triple::Hull:
971 case llvm::Triple::Domain:
972 case llvm::Triple::RayGeneration:
973 case llvm::Triple::Intersection:
974 case llvm::Triple::AnyHit:
975 case llvm::Triple::ClosestHit:
976 case llvm::Triple::Miss:
977 case llvm::Triple::Callable:
978 if (
const auto *NT = FD->
getAttr<HLSLNumThreadsAttr>()) {
979 diagnoseAttrStageMismatch(NT, ST,
980 {llvm::Triple::Compute,
981 llvm::Triple::Amplification,
982 llvm::Triple::Mesh});
985 if (
const auto *WS = FD->
getAttr<HLSLWaveSizeAttr>()) {
986 diagnoseAttrStageMismatch(WS, ST,
987 {llvm::Triple::Compute,
988 llvm::Triple::Amplification,
989 llvm::Triple::Mesh});
994 case llvm::Triple::Compute:
995 case llvm::Triple::Amplification:
996 case llvm::Triple::Mesh:
997 if (!FD->
hasAttr<HLSLNumThreadsAttr>()) {
999 << llvm::Triple::getEnvironmentTypeName(ST);
1002 if (
const auto *WS = FD->
getAttr<HLSLWaveSizeAttr>()) {
1003 if (Ver < VersionTuple(6, 6)) {
1004 Diag(WS->getLocation(), diag::err_hlsl_attribute_in_wrong_shader_model)
1007 }
else if (WS->getSpelledArgsCount() > 1 && Ver < VersionTuple(6, 8)) {
1010 diag::err_hlsl_attribute_number_arguments_insufficient_shader_model)
1011 << WS << WS->getSpelledArgsCount() <<
"6.8";
1016 case llvm::Triple::RootSignature:
1017 llvm_unreachable(
"rootsig environment has no function entry point");
1019 llvm_unreachable(
"Unhandled environment in triple");
1022 SemaHLSL::SemanticContext InputSC = {};
1023 InputSC.CurrentIOType = IOType::In;
1026 SemanticInfo ActiveSemantic;
1027 ActiveSemantic.Semantic = Param->getAttr<HLSLParsedSemanticAttr>();
1028 if (ActiveSemantic.Semantic)
1029 ActiveSemantic.Index = ActiveSemantic.Semantic->getSemanticIndex();
1032 if (!determineActiveSemantic(FD, Param, Param, ActiveSemantic, InputSC)) {
1033 Diag(Param->getLocation(), diag::note_previous_decl) << Param;
1038 SemanticInfo ActiveSemantic;
1039 SemaHLSL::SemanticContext OutputSC = {};
1040 OutputSC.CurrentIOType = IOType::Out;
1041 ActiveSemantic.Semantic = FD->
getAttr<HLSLParsedSemanticAttr>();
1042 if (ActiveSemantic.Semantic)
1043 ActiveSemantic.Index = ActiveSemantic.Semantic->getSemanticIndex();
1045 determineActiveSemantic(FD, FD, FD, ActiveSemantic, OutputSC);
1048void SemaHLSL::checkSemanticAnnotation(
1050 const HLSLAppliedSemanticAttr *SemanticAttr,
const SemanticContext &SC) {
1051 auto *ShaderAttr = EntryPoint->
getAttr<HLSLShaderAttr>();
1052 assert(ShaderAttr &&
"Entry point has no shader attribute");
1053 llvm::Triple::EnvironmentType ST = ShaderAttr->getType();
1055 auto SemanticName = SemanticAttr->getSemanticName().upper();
1056 if (SemanticName ==
"SV_DISPATCHTHREADID" ||
1057 SemanticName ==
"SV_GROUPINDEX" || SemanticName ==
"SV_GROUPTHREADID" ||
1058 SemanticName ==
"SV_GROUPID") {
1060 if (ST != llvm::Triple::Compute)
1061 diagnoseSemanticStageMismatch(SemanticAttr, ST, SC.CurrentIOType,
1062 {{llvm::Triple::Compute, IOType::In}});
1064 if (SemanticAttr->getSemanticIndex() != 0) {
1065 std::string PrettyName =
1066 "'" + SemanticAttr->getSemanticName().str() +
"'";
1067 Diag(SemanticAttr->getLoc(),
1068 diag::err_hlsl_semantic_indexing_not_supported)
1074 if (SemanticName ==
"SV_POSITION") {
1077 diagnoseSemanticStageMismatch(SemanticAttr, ST, SC.CurrentIOType,
1078 {{llvm::Triple::Vertex, IOType::InOut},
1079 {llvm::Triple::Pixel, IOType::In}});
1082 if (SemanticName ==
"SV_VERTEXID") {
1083 diagnoseSemanticStageMismatch(SemanticAttr, ST, SC.CurrentIOType,
1084 {{llvm::Triple::Vertex, IOType::In}});
1088 if (SemanticName ==
"SV_TARGET") {
1089 diagnoseSemanticStageMismatch(SemanticAttr, ST, SC.CurrentIOType,
1090 {{llvm::Triple::Pixel, IOType::Out}});
1096 if (SemanticAttr->getAttrName()->getName().starts_with_insensitive(
"SV_"))
1097 llvm_unreachable(
"Unknown SemanticAttr");
1100void SemaHLSL::diagnoseAttrStageMismatch(
1101 const Attr *A, llvm::Triple::EnvironmentType Stage,
1102 std::initializer_list<llvm::Triple::EnvironmentType> AllowedStages) {
1103 SmallVector<StringRef, 8> StageStrings;
1104 llvm::transform(AllowedStages, std::back_inserter(StageStrings),
1105 [](llvm::Triple::EnvironmentType ST) {
1107 HLSLShaderAttr::ConvertEnvironmentTypeToStr(ST));
1109 Diag(A->
getLoc(), diag::err_hlsl_attr_unsupported_in_stage)
1110 << A->
getAttrName() << llvm::Triple::getEnvironmentTypeName(Stage)
1111 << (AllowedStages.size() != 1) << join(StageStrings,
", ");
1114void SemaHLSL::diagnoseSemanticStageMismatch(
1115 const Attr *A, llvm::Triple::EnvironmentType Stage, IOType CurrentIOType,
1116 std::initializer_list<SemanticStageInfo> Allowed) {
1118 for (
auto &Case : Allowed) {
1119 if (Case.Stage != Stage)
1122 if (CurrentIOType & Case.AllowedIOTypesMask)
1125 SmallVector<std::string, 8> ValidCases;
1127 Allowed, std::back_inserter(ValidCases), [](SemanticStageInfo Case) {
1128 SmallVector<std::string, 2> ValidType;
1129 if (Case.AllowedIOTypesMask & IOType::In)
1130 ValidType.push_back(
"input");
1131 if (Case.AllowedIOTypesMask & IOType::Out)
1132 ValidType.push_back(
"output");
1134 HLSLShaderAttr::ConvertEnvironmentTypeToStr(Case.Stage)) +
1135 " " + join(ValidType,
"/");
1137 Diag(A->
getLoc(), diag::err_hlsl_semantic_unsupported_iotype_for_stage)
1138 << A->
getAttrName() << (CurrentIOType & IOType::In ?
"input" :
"output")
1139 << llvm::Triple::getEnvironmentTypeName(Case.Stage)
1140 << join(ValidCases,
", ");
1144 SmallVector<StringRef, 8> StageStrings;
1146 Allowed, std::back_inserter(StageStrings), [](SemanticStageInfo Case) {
1148 HLSLShaderAttr::ConvertEnvironmentTypeToStr(Case.Stage));
1151 Diag(A->
getLoc(), diag::err_hlsl_attr_unsupported_in_stage)
1152 << A->
getAttrName() << llvm::Triple::getEnvironmentTypeName(Stage)
1153 << (Allowed.size() != 1) << join(StageStrings,
", ");
1156template <CastKind Kind>
1159 Ty = VTy->getElementType();
1164template <CastKind Kind>
1176 if (LHSFloat && RHSFloat) {
1204 if (LHSSigned == RHSSigned) {
1205 if (IsCompAssign || IntOrder >= 0)
1213 if (IntOrder != (LHSSigned ? 1 : -1)) {
1214 if (IsCompAssign || RHSSigned)
1222 if (Ctx.getIntWidth(LElTy) != Ctx.getIntWidth(RElTy)) {
1223 if (IsCompAssign || LHSSigned)
1239 QualType ElTy = Ctx.getCorrespondingUnsignedType(LHSSigned ? LElTy : RElTy);
1240 QualType NewTy = Ctx.getExtVectorType(
1250 return CK_FloatingCast;
1252 return CK_IntegralCast;
1254 return CK_IntegralToFloating;
1256 return CK_FloatingToIntegral;
1262 bool IsCompAssign) {
1269 if (!LVecTy && IsCompAssign) {
1271 RHS =
SemaRef.ImpCastExprToType(RHS.
get(), RElTy, CK_HLSLVectorTruncation);
1273 if (Ctx.hasSameUnqualifiedType(LHSType, RHSType))
1275 RHS =
SemaRef.ImpCastExprToType(RHS.
get(), LHSType,
1280 unsigned EndSz = std::numeric_limits<unsigned>::max();
1283 LSz = EndSz = LVecTy->getNumElements();
1286 assert(EndSz != std::numeric_limits<unsigned>::max() &&
1287 "one of the above should have had a value");
1291 if (IsCompAssign && LSz != EndSz) {
1293 diag::err_hlsl_vector_compound_assignment_truncation)
1294 << LHSType << RHSType;
1300 if (!IsCompAssign && LVecTy && LVecTy->getNumElements() > EndSz)
1305 if (!IsCompAssign && !LVecTy)
1309 if (Ctx.hasSameUnqualifiedType(LHSType, RHSType))
1310 return Ctx.getCommonSugaredType(LHSType, RHSType);
1318 LElTy, RElTy, IsCompAssign);
1321 "HLSL Vectors can only contain integer or floating point types");
1323 LElTy, RElTy, IsCompAssign);
1328 assert((Opc == BO_LOr || Opc == BO_LAnd) &&
1329 "Called with non-logical operator");
1331 llvm::raw_svector_ostream OS(Buff);
1333 StringRef NewFnName = Opc == BO_LOr ?
"or" :
"and";
1334 OS << NewFnName <<
"(";
1344std::pair<IdentifierInfo *, bool>
1347 std::string IdStr =
"__hlsl_rootsig_decl_" + std::to_string(Hash);
1354 return {DeclIdent,
Found};
1365 for (
auto &RootSigElement : RootElements)
1366 Elements.push_back(RootSigElement.getElement());
1370 DeclIdent,
SemaRef.getLangOpts().HLSLRootSigVer, Elements);
1372 SignatureDecl->setImplicit();
1378 if (RootSigOverrideIdent) {
1381 if (
SemaRef.LookupQualifiedName(R, DC))
1382 return dyn_cast<HLSLRootSignatureDecl>(R.getFoundDecl());
1390struct PerVisibilityBindingChecker {
1393 std::array<llvm::hlsl::BindingInfoBuilder, 8> Builders;
1397 llvm::dxbc::ShaderVisibility Vis;
1402 PerVisibilityBindingChecker(
SemaHLSL *S) : S(S) {}
1404 void trackBinding(llvm::dxbc::ShaderVisibility
Visibility,
1405 llvm::dxil::ResourceClass RC, uint32_t Space,
1406 uint32_t LowerBound, uint32_t UpperBound,
1407 const hlsl::RootSignatureElement *Elem) {
1409 assert(BuilderIndex < Builders.size() &&
1410 "Not enough builders for visibility type");
1411 Builders[BuilderIndex].trackBinding(RC, Space, LowerBound, UpperBound,
1412 static_cast<const void *
>(Elem));
1414 static_assert(llvm::to_underlying(llvm::dxbc::ShaderVisibility::All) == 0,
1415 "'All' visibility must come first");
1416 if (
Visibility == llvm::dxbc::ShaderVisibility::All)
1417 for (
size_t I = 1, E = Builders.size(); I < E; ++I)
1418 Builders[I].trackBinding(RC, Space, LowerBound, UpperBound,
1419 static_cast<const void *
>(Elem));
1421 ElemInfoMap.push_back({Elem,
Visibility,
false});
1424 ElemInfo &
getInfo(
const hlsl::RootSignatureElement *Elem) {
1425 auto It = llvm::lower_bound(
1427 [](
const auto &LHS,
const auto &RHS) {
return LHS.Elem < RHS; });
1428 assert(It->Elem == Elem &&
"Element not in map");
1432 bool checkOverlap() {
1433 llvm::sort(ElemInfoMap, [](
const auto &LHS,
const auto &RHS) {
1434 return LHS.Elem < RHS.Elem;
1437 bool HadOverlap =
false;
1439 using llvm::hlsl::BindingInfoBuilder;
1440 auto ReportOverlap = [
this,
1441 &HadOverlap](
const BindingInfoBuilder &Builder,
1442 const llvm::hlsl::Binding &Reported) {
1446 static_cast<const hlsl::RootSignatureElement *
>(Reported.Cookie);
1447 const llvm::hlsl::Binding &
Previous = Builder.findOverlapping(Reported);
1448 const auto *PrevElem =
1449 static_cast<const hlsl::RootSignatureElement *
>(
Previous.Cookie);
1451 ElemInfo &Info =
getInfo(Elem);
1456 Info.Diagnosed =
true;
1458 ElemInfo &PrevInfo =
getInfo(PrevElem);
1459 llvm::dxbc::ShaderVisibility CommonVis =
1460 Info.Vis == llvm::dxbc::ShaderVisibility::All ? PrevInfo.Vis
1463 this->S->
Diag(Elem->
getLocation(), diag::err_hlsl_resource_range_overlap)
1464 << llvm::to_underlying(Reported.RC) << Reported.LowerBound
1465 << Reported.isUnbounded() << Reported.UpperBound
1470 this->S->
Diag(PrevElem->getLocation(),
1471 diag::note_hlsl_resource_range_here);
1474 for (BindingInfoBuilder &Builder : Builders)
1475 Builder.calculateBindingInfo(ReportOverlap);
1499 if (
const auto *ResTy =
1500 SecondField->
getType()->
getAs<HLSLAttributedResourceType>()) {
1501 return ResTy->getAttrs().IsCounter;
1509 bool HadError =
false;
1510 auto ReportError = [
this, &HadError](
SourceLocation Loc, uint32_t LowerBound,
1511 uint32_t UpperBound) {
1513 this->
Diag(Loc, diag::err_hlsl_invalid_rootsig_value)
1514 << LowerBound << UpperBound;
1521 this->
Diag(Loc, diag::err_hlsl_invalid_rootsig_value)
1522 << llvm::formatv(
"{0:f}", LowerBound).sstr<6>()
1523 << llvm::formatv(
"{0:f}", UpperBound).sstr<6>();
1526 auto VerifyRegister = [ReportError](
SourceLocation Loc, uint32_t Register) {
1527 if (!llvm::hlsl::rootsig::verifyRegisterValue(Register))
1528 ReportError(Loc, 0, 0xfffffffe);
1531 auto VerifySpace = [ReportError](
SourceLocation Loc, uint32_t Space) {
1532 if (!llvm::hlsl::rootsig::verifyRegisterSpace(Space))
1533 ReportError(Loc, 0, 0xffffffef);
1536 const uint32_t Version =
1537 llvm::to_underlying(
SemaRef.getLangOpts().HLSLRootSigVer);
1538 const uint32_t VersionEnum = Version - 1;
1539 auto ReportFlagError = [
this, &HadError, VersionEnum](
SourceLocation Loc) {
1541 this->
Diag(Loc, diag::err_hlsl_invalid_rootsig_flag)
1548 const llvm::hlsl::rootsig::RootElement &Elem = RootSigElem.
getElement();
1549 if (
const auto *Descriptor =
1550 std::get_if<llvm::hlsl::rootsig::RootDescriptor>(&Elem)) {
1551 VerifyRegister(Loc, Descriptor->Reg.Number);
1552 VerifySpace(Loc, Descriptor->Space);
1554 if (!llvm::hlsl::rootsig::verifyRootDescriptorFlag(Version,
1556 ReportFlagError(Loc);
1557 }
else if (
const auto *Constants =
1558 std::get_if<llvm::hlsl::rootsig::RootConstants>(&Elem)) {
1559 VerifyRegister(Loc, Constants->Reg.Number);
1560 VerifySpace(Loc, Constants->Space);
1561 }
else if (
const auto *Sampler =
1562 std::get_if<llvm::hlsl::rootsig::StaticSampler>(&Elem)) {
1563 VerifyRegister(Loc, Sampler->Reg.Number);
1564 VerifySpace(Loc, Sampler->Space);
1567 "By construction, parseFloatParam can't produce a NaN from a "
1568 "float_literal token");
1570 if (!llvm::hlsl::rootsig::verifyMaxAnisotropy(Sampler->MaxAnisotropy))
1571 ReportError(Loc, 0, 16);
1572 if (!llvm::hlsl::rootsig::verifyMipLODBias(Sampler->MipLODBias))
1573 ReportFloatError(Loc, -16.f, 15.99f);
1574 }
else if (
const auto *Clause =
1575 std::get_if<llvm::hlsl::rootsig::DescriptorTableClause>(
1577 VerifyRegister(Loc, Clause->Reg.Number);
1578 VerifySpace(Loc, Clause->Space);
1580 if (!llvm::hlsl::rootsig::verifyNumDescriptors(Clause->NumDescriptors)) {
1584 ReportError(Loc, 1, 0xfffffffe);
1587 if (!llvm::hlsl::rootsig::verifyDescriptorRangeFlag(Version, Clause->Type,
1589 ReportFlagError(Loc);
1593 PerVisibilityBindingChecker BindingChecker(
this);
1594 SmallVector<std::pair<
const llvm::hlsl::rootsig::DescriptorTableClause *,
1599 const llvm::hlsl::rootsig::RootElement &Elem = RootSigElem.
getElement();
1600 if (
const auto *Descriptor =
1601 std::get_if<llvm::hlsl::rootsig::RootDescriptor>(&Elem)) {
1602 uint32_t LowerBound(Descriptor->Reg.Number);
1603 uint32_t UpperBound(LowerBound);
1605 BindingChecker.trackBinding(
1606 Descriptor->Visibility,
1607 static_cast<llvm::dxil::ResourceClass
>(Descriptor->Type),
1608 Descriptor->Space, LowerBound, UpperBound, &RootSigElem);
1609 }
else if (
const auto *Constants =
1610 std::get_if<llvm::hlsl::rootsig::RootConstants>(&Elem)) {
1611 uint32_t LowerBound(Constants->Reg.Number);
1612 uint32_t UpperBound(LowerBound);
1614 BindingChecker.trackBinding(
1615 Constants->Visibility, llvm::dxil::ResourceClass::CBuffer,
1616 Constants->Space, LowerBound, UpperBound, &RootSigElem);
1617 }
else if (
const auto *Sampler =
1618 std::get_if<llvm::hlsl::rootsig::StaticSampler>(&Elem)) {
1619 uint32_t LowerBound(Sampler->Reg.Number);
1620 uint32_t UpperBound(LowerBound);
1622 BindingChecker.trackBinding(
1623 Sampler->Visibility, llvm::dxil::ResourceClass::Sampler,
1624 Sampler->Space, LowerBound, UpperBound, &RootSigElem);
1625 }
else if (
const auto *Clause =
1626 std::get_if<llvm::hlsl::rootsig::DescriptorTableClause>(
1629 UnboundClauses.emplace_back(Clause, &RootSigElem);
1630 }
else if (
const auto *Table =
1631 std::get_if<llvm::hlsl::rootsig::DescriptorTable>(&Elem)) {
1632 assert(UnboundClauses.size() == Table->NumClauses &&
1633 "Number of unbound elements must match the number of clauses");
1634 bool HasAnySampler =
false;
1635 bool HasAnyNonSampler =
false;
1636 uint64_t Offset = 0;
1637 bool IsPrevUnbound =
false;
1638 for (
const auto &[Clause, ClauseElem] : UnboundClauses) {
1640 if (Clause->Type == llvm::dxil::ResourceClass::Sampler)
1641 HasAnySampler =
true;
1643 HasAnyNonSampler =
true;
1645 if (HasAnySampler && HasAnyNonSampler)
1646 Diag(Loc, diag::err_hlsl_invalid_mixed_resources);
1651 if (Clause->NumDescriptors == 0)
1655 Clause->Offset == llvm::hlsl::rootsig::DescriptorTableOffsetAppend;
1657 Offset = Clause->Offset;
1659 uint64_t RangeBound = llvm::hlsl::rootsig::computeRangeBound(
1660 Offset, Clause->NumDescriptors);
1662 if (IsPrevUnbound && IsAppending)
1663 Diag(Loc, diag::err_hlsl_appending_onto_unbound);
1664 else if (!llvm::hlsl::rootsig::verifyNoOverflowedOffset(RangeBound))
1665 Diag(Loc, diag::err_hlsl_offset_overflow) << Offset << RangeBound;
1668 Offset = RangeBound + 1;
1669 IsPrevUnbound = Clause->NumDescriptors ==
1670 llvm::hlsl::rootsig::NumDescriptorsUnbounded;
1673 uint32_t LowerBound(Clause->Reg.Number);
1674 uint32_t UpperBound = llvm::hlsl::rootsig::computeRangeBound(
1675 LowerBound, Clause->NumDescriptors);
1677 BindingChecker.trackBinding(
1679 static_cast<llvm::dxil::ResourceClass
>(Clause->Type), Clause->Space,
1680 LowerBound, UpperBound, ClauseElem);
1682 UnboundClauses.clear();
1686 return BindingChecker.checkOverlap();
1691 Diag(AL.
getLoc(), diag::err_attribute_wrong_number_arguments) << AL << 1;
1696 if (
auto *RS = D->
getAttr<RootSignatureAttr>()) {
1697 if (RS->getSignatureIdent() != Ident) {
1698 Diag(AL.
getLoc(), diag::err_disallowed_duplicate_attribute) << RS;
1702 Diag(AL.
getLoc(), diag::warn_duplicate_attribute_exact) << RS;
1708 if (
auto *SignatureDecl =
1709 dyn_cast<HLSLRootSignatureDecl>(R.getFoundDecl())) {
1716 llvm::VersionTuple SMVersion =
1721 uint32_t ZMax = 1024;
1722 uint32_t ThreadMax = 1024;
1723 if (IsDXIL && SMVersion.getMajor() <= 4) {
1726 }
else if (IsDXIL && SMVersion.getMajor() == 5) {
1736 diag::err_hlsl_numthreads_argument_oor)
1745 diag::err_hlsl_numthreads_argument_oor)
1754 diag::err_hlsl_numthreads_argument_oor)
1759 if (
X * Y * Z > ThreadMax) {
1760 Diag(AL.
getLoc(), diag::err_hlsl_numthreads_invalid) << ThreadMax;
1777 if (SpelledArgsCount == 0 || SpelledArgsCount > 3)
1785 if (SpelledArgsCount > 1 &&
1789 uint32_t Preferred = 0;
1790 if (SpelledArgsCount > 2 &&
1794 if (SpelledArgsCount > 2) {
1797 diag::err_attribute_power_of_two_in_range)
1798 << AL << llvm::dxil::MinWaveSize << llvm::dxil::MaxWaveSize
1803 if (Preferred < Min || Preferred >
Max) {
1805 diag::err_attribute_power_of_two_in_range)
1806 << AL <<
Min <<
Max << Preferred;
1809 }
else if (SpelledArgsCount > 1) {
1812 diag::err_attribute_power_of_two_in_range)
1813 << AL << llvm::dxil::MinWaveSize << llvm::dxil::MaxWaveSize <<
Max;
1817 Diag(AL.
getLoc(), diag::err_attribute_argument_invalid) << AL << 1;
1820 Diag(AL.
getLoc(), diag::warn_attr_min_eq_max) << AL;
1825 diag::err_attribute_power_of_two_in_range)
1826 << AL << llvm::dxil::MinWaveSize << llvm::dxil::MaxWaveSize <<
Min;
1831 HLSLWaveSizeAttr *NewAttr =
1868 uint32_t Binding = 0;
1892 if (!T->hasUnsignedIntegerRepresentation() ||
1893 (VT && VT->getNumElements() > 3)) {
1894 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type)
1895 << AL <<
"uint/uint2/uint3";
1904 if (!T->hasFloatingRepresentation() || (VT && VT->getNumElements() > 4)) {
1905 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type)
1906 << AL <<
"float/float1/float2/float3/float4";
1914 std::optional<unsigned> Index) {
1918 QualType ValueType = VD->getType();
1919 if (
auto *FD = dyn_cast<FunctionDecl>(D))
1922 bool IsOutput =
false;
1923 if (HLSLParamModifierAttr *MA = D->
getAttr<HLSLParamModifierAttr>()) {
1930 if (SemanticName ==
"SV_DISPATCHTHREADID") {
1933 Diag(AL.
getLoc(), diag::err_hlsl_semantic_output_not_supported) << AL;
1934 if (Index.has_value())
1935 Diag(AL.
getLoc(), diag::err_hlsl_semantic_indexing_not_supported) << AL;
1940 if (SemanticName ==
"SV_GROUPINDEX") {
1942 Diag(AL.
getLoc(), diag::err_hlsl_semantic_output_not_supported) << AL;
1943 if (Index.has_value())
1944 Diag(AL.
getLoc(), diag::err_hlsl_semantic_indexing_not_supported) << AL;
1949 if (SemanticName ==
"SV_GROUPTHREADID") {
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_GROUPID") {
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_POSITION") {
1970 const auto *VT = ValueType->getAs<
VectorType>();
1971 if (!ValueType->hasFloatingRepresentation() ||
1972 (VT && VT->getNumElements() > 4))
1973 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type)
1974 << AL <<
"float/float1/float2/float3/float4";
1979 if (SemanticName ==
"SV_VERTEXID") {
1980 uint64_t SizeInBits =
SemaRef.Context.getTypeSize(ValueType);
1981 if (!ValueType->isUnsignedIntegerType() || SizeInBits != 32)
1982 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type) << AL <<
"uint";
1987 if (SemanticName ==
"SV_TARGET") {
1988 const auto *VT = ValueType->getAs<
VectorType>();
1989 if (!ValueType->hasFloatingRepresentation() ||
1990 (VT && VT->getNumElements() > 4))
1991 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type)
1992 << AL <<
"float/float1/float2/float3/float4";
1997 Diag(AL.
getLoc(), diag::err_hlsl_unknown_semantic) << AL;
2001 uint32_t IndexValue(0), ExplicitIndex(0);
2004 assert(0 &&
"HLSLUnparsedSemantic is expected to have 2 int arguments.");
2006 assert(IndexValue > 0 ? ExplicitIndex :
true);
2007 std::optional<unsigned> Index =
2008 ExplicitIndex ? std::optional<unsigned>(IndexValue) : std::nullopt;
2018 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_ast_node)
2019 << AL <<
"shader constant in a constant buffer";
2023 uint32_t SubComponent;
2033 bool IsAggregateTy = (T->isArrayType() || T->isStructureType());
2038 if (IsAggregateTy) {
2039 Diag(AL.
getLoc(), diag::err_hlsl_invalid_register_or_packoffset);
2043 if ((Component * 32 + Size) > 128) {
2044 Diag(AL.
getLoc(), diag::err_hlsl_packoffset_cross_reg_boundary);
2049 EltTy = VT->getElementType();
2051 if (Align > 32 && Component == 1) {
2054 Diag(AL.
getLoc(), diag::err_hlsl_packoffset_alignment_mismatch)
2068 if (!
SemaRef.checkStringLiteralArgumentAttr(AL, 0, Str, &ArgLoc))
2071 llvm::Triple::EnvironmentType ShaderType;
2072 if (!HLSLShaderAttr::ConvertStrToEnvironmentType(Str, ShaderType)) {
2073 Diag(AL.
getLoc(), diag::warn_attribute_type_not_supported)
2074 << AL << Str << ArgLoc;
2088 assert(AttrList.size() &&
"expected list of resource attributes");
2095 HLSLAttributedResourceType::Attributes ResAttrs;
2097 bool HasResourceClass =
false;
2098 bool HasResourceDimension =
false;
2099 for (
const Attr *A : AttrList) {
2104 case attr::HLSLResourceClass: {
2106 if (HasResourceClass) {
2108 ? diag::warn_duplicate_attribute_exact
2109 : diag::warn_duplicate_attribute)
2113 ResAttrs.ResourceClass = RC;
2114 HasResourceClass =
true;
2117 case attr::HLSLResourceDimension: {
2118 llvm::dxil::ResourceDimension RD =
2120 if (HasResourceDimension) {
2122 ? diag::warn_duplicate_attribute_exact
2123 : diag::warn_duplicate_attribute)
2127 ResAttrs.ResourceDimension = RD;
2128 HasResourceDimension =
true;
2132 if (ResAttrs.IsROV) {
2136 ResAttrs.IsROV =
true;
2138 case attr::HLSLRawBuffer:
2139 if (ResAttrs.RawBuffer) {
2143 ResAttrs.RawBuffer =
true;
2145 case attr::HLSLIsCounter:
2146 if (ResAttrs.IsCounter) {
2150 ResAttrs.IsCounter =
true;
2152 case attr::HLSLContainedType: {
2155 if (!ContainedTy.
isNull()) {
2157 ? diag::warn_duplicate_attribute_exact
2158 : diag::warn_duplicate_attribute)
2167 llvm_unreachable(
"unhandled resource attribute type");
2171 if (!HasResourceClass) {
2172 S.
Diag(AttrList.back()->getRange().getEnd(),
2173 diag::err_hlsl_missing_resource_class);
2178 Wrapped, ContainedTy, ResAttrs);
2180 if (LocInfo && ContainedTyInfo) {
2193 if (!T->isHLSLResourceType()) {
2194 Diag(AL.
getLoc(), diag::err_hlsl_attribute_needs_intangible_type)
2209 AttributeCommonInfo::AS_CXX11, 0, false ,
2214 case ParsedAttr::AT_HLSLResourceClass: {
2216 Diag(AL.
getLoc(), diag::err_attribute_argument_type)
2227 if (!HLSLResourceClassAttr::ConvertStrToResourceClass(Identifier, RC)) {
2228 Diag(ArgLoc, diag::warn_attribute_type_not_supported)
2229 <<
"ResourceClass" << Identifier;
2232 A = HLSLResourceClassAttr::Create(
getASTContext(), RC, ACI);
2236 case ParsedAttr::AT_HLSLResourceDimension: {
2237 StringRef Identifier;
2239 if (!
SemaRef.checkStringLiteralArgumentAttr(AL, 0, Identifier, &ArgLoc))
2243 llvm::dxil::ResourceDimension RD;
2244 if (!HLSLResourceDimensionAttr::ConvertStrToResourceDimension(Identifier,
2246 Diag(ArgLoc, diag::warn_attribute_type_not_supported)
2247 <<
"ResourceDimension" << Identifier;
2250 A = HLSLResourceDimensionAttr::Create(
getASTContext(), RD, ACI);
2254 case ParsedAttr::AT_HLSLROV:
2258 case ParsedAttr::AT_HLSLRawBuffer:
2262 case ParsedAttr::AT_HLSLIsCounter:
2266 case ParsedAttr::AT_HLSLContainedType: {
2268 Diag(AL.
getLoc(), diag::err_attribute_wrong_number_arguments) << AL << 1;
2274 assert(TSI &&
"no type source info for attribute argument");
2276 diag::err_incomplete_type))
2278 A = HLSLContainedTypeAttr::Create(
getASTContext(), TSI, ACI);
2283 llvm_unreachable(
"unhandled HLSL attribute");
2286 HLSLResourcesTypeAttrs.emplace_back(A);
2292 if (!HLSLResourcesTypeAttrs.size())
2298 HLSLResourcesTypeAttrs, QT, &LocInfo)) {
2299 const HLSLAttributedResourceType *RT =
2306 LocsForHLSLAttributedResources.insert(std::pair(RT, LocInfo));
2308 HLSLResourcesTypeAttrs.clear();
2316 auto I = LocsForHLSLAttributedResources.find(RT);
2317 if (I != LocsForHLSLAttributedResources.end()) {
2318 LocInfo = I->second;
2319 LocsForHLSLAttributedResources.erase(I);
2328void SemaHLSL::collectResourceBindingsOnUserRecordDecl(
const VarDecl *VD,
2329 const RecordType *RT) {
2330 const RecordDecl *RD = RT->getDecl()->getDefinitionOrSelf();
2337 "incomplete arrays inside user defined types are not supported");
2346 if (
const HLSLAttributedResourceType *AttrResType =
2347 HLSLAttributedResourceType::findHandleTypeOnResource(Ty)) {
2352 Bindings.addDeclBindingInfo(VD, RC);
2353 }
else if (
const RecordType *RT = dyn_cast<RecordType>(Ty)) {
2359 collectResourceBindingsOnUserRecordDecl(VD, RT);
2371 bool SpecifiedSpace) {
2372 int RegTypeNum =
static_cast<int>(RegType);
2375 if (D->
hasAttr<HLSLGroupSharedAddressSpaceAttr>()) {
2376 S.
Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
2381 if (
HLSLBufferDecl *CBufferOrTBuffer = dyn_cast<HLSLBufferDecl>(D)) {
2382 ResourceClass RC = CBufferOrTBuffer->isCBuffer() ? ResourceClass::CBuffer
2383 : ResourceClass::SRV;
2393 assert(
isa<VarDecl>(D) &&
"D is expected to be VarDecl or HLSLBufferDecl");
2397 if (
const HLSLAttributedResourceType *AttrResType =
2398 HLSLAttributedResourceType::findHandleTypeOnResource(
2415 if (SpecifiedSpace && !DeclaredInCOrTBuffer)
2416 S.
Diag(ArgLoc, diag::err_hlsl_space_on_global_constant);
2421 if (RegType == RegisterType::CBuffer)
2422 S.
Diag(ArgLoc, diag::warn_hlsl_deprecated_register_type_b);
2423 else if (RegType != RegisterType::C)
2424 S.
Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
2428 if (RegType == RegisterType::C)
2429 S.
Diag(ArgLoc, diag::warn_hlsl_register_type_c_packoffset);
2431 S.
Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
2441 S.
Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
2449 bool RegisterTypesDetected[5] = {
false};
2450 RegisterTypesDetected[
static_cast<int>(regType)] =
true;
2453 if (HLSLResourceBindingAttr *
attr =
2454 dyn_cast<HLSLResourceBindingAttr>(*it)) {
2457 if (RegisterTypesDetected[
static_cast<int>(otherRegType)]) {
2458 int otherRegTypeNum =
static_cast<int>(otherRegType);
2460 diag::err_hlsl_duplicate_register_annotation)
2464 RegisterTypesDetected[
static_cast<int>(otherRegType)] =
true;
2472 bool SpecifiedSpace) {
2477 "expecting VarDecl or HLSLBufferDecl");
2489 const uint64_t &Limit,
2492 uint64_t ArrayCount = 1) {
2497 if (StartSlot > Limit)
2501 if (
const auto *AT = dyn_cast<ArrayType>(T)) {
2504 if (
const auto *CAT = dyn_cast<ConstantArrayType>(AT))
2505 Count = CAT->
getSize().getZExtValue();
2509 ArrayCount * Count);
2513 if (
auto ResTy = dyn_cast<HLSLAttributedResourceType>(T)) {
2516 if (ResTy->getAttrs().ResourceClass != ResClass)
2520 uint64_t EndSlot = StartSlot + ArrayCount - 1;
2521 if (EndSlot > Limit)
2525 StartSlot = EndSlot + 1;
2530 if (
const auto *RT = dyn_cast<RecordType>(T)) {
2533 if (
const auto *CXXRD = dyn_cast<CXXRecordDecl>(RD)) {
2536 ResClass, Ctx, ArrayCount))
2543 ResClass, Ctx, ArrayCount))
2557 const uint64_t Limit = UINT32_MAX;
2558 if (SlotNum > Limit)
2563 if (RegTy == RegisterType::C || RegTy == RegisterType::I)
2566 if (
VarDecl *VD = dyn_cast<VarDecl>(TheDecl)) {
2567 uint64_t BaseSlot = SlotNum;
2575 return (BaseSlot > Limit);
2582 return (SlotNum > Limit);
2585 llvm_unreachable(
"unexpected decl type");
2589 if (
VarDecl *VD = dyn_cast<VarDecl>(TheDecl)) {
2591 if (
const auto *IAT = dyn_cast<IncompleteArrayType>(Ty))
2592 Ty = IAT->getElementType();
2594 diag::err_incomplete_type))
2598 StringRef Slot =
"";
2599 StringRef Space =
"";
2603 Diag(AL.
getLoc(), diag::err_attribute_argument_type)
2613 Diag(AL.
getLoc(), diag::err_attribute_argument_type)
2619 SpaceLoc = Loc->
getLoc();
2622 if (Str.starts_with(
"space")) {
2624 SpaceLoc = Loc->
getLoc();
2633 std::optional<unsigned> SlotNum;
2634 unsigned SpaceNum = 0;
2637 if (!Slot.empty()) {
2639 Diag(SlotLoc, diag::err_hlsl_binding_type_invalid) << Slot.substr(0, 1);
2642 if (RegType == RegisterType::I) {
2643 Diag(SlotLoc, diag::warn_hlsl_deprecated_register_type_i);
2646 const StringRef SlotNumStr = Slot.substr(1);
2651 if (SlotNumStr.getAsInteger(10, N)) {
2652 Diag(SlotLoc, diag::err_hlsl_unsupported_register_number);
2660 Diag(SlotLoc, diag::err_hlsl_register_number_too_large);
2669 if (!Space.starts_with(
"space")) {
2670 Diag(SpaceLoc, diag::err_hlsl_expected_space) << Space;
2673 StringRef SpaceNumStr = Space.substr(5);
2674 if (SpaceNumStr.getAsInteger(10, SpaceNum)) {
2675 Diag(SpaceLoc, diag::err_hlsl_expected_space) << Space;
2680 if (SlotNum.has_value())
2685 HLSLResourceBindingAttr *NewAttr =
2686 HLSLResourceBindingAttr::Create(
getASTContext(), Slot, Space, AL);
2688 NewAttr->setBinding(RegType, SlotNum, SpaceNum);
2743 llvm::DenseMap<const FunctionDecl *, unsigned> ScannedDecls;
2747 llvm::Triple::EnvironmentType CurrentShaderEnvironment;
2748 unsigned CurrentShaderStageBit;
2753 bool ReportOnlyShaderStageIssues;
2756 void SetShaderStageContext(llvm::Triple::EnvironmentType ShaderType) {
2757 static_assert(
sizeof(
unsigned) >= 4);
2758 assert(HLSLShaderAttr::isValidShaderType(ShaderType));
2759 assert((
unsigned)(ShaderType - llvm::Triple::Pixel) < 31 &&
2760 "ShaderType is too big for this bitmap");
2763 unsigned bitmapIndex = ShaderType - llvm::Triple::Pixel;
2764 CurrentShaderEnvironment = ShaderType;
2765 CurrentShaderStageBit = (1 << bitmapIndex);
2768 void SetUnknownShaderStageContext() {
2769 CurrentShaderEnvironment = llvm::Triple::UnknownEnvironment;
2770 CurrentShaderStageBit = (1 << 31);
2773 llvm::Triple::EnvironmentType GetCurrentShaderEnvironment()
const {
2774 return CurrentShaderEnvironment;
2777 bool InUnknownShaderStageContext()
const {
2778 return CurrentShaderEnvironment == llvm::Triple::UnknownEnvironment;
2782 void AddToScannedFunctions(
const FunctionDecl *FD) {
2783 unsigned &ScannedStages = ScannedDecls[FD];
2784 ScannedStages |= CurrentShaderStageBit;
2787 unsigned GetScannedStages(
const FunctionDecl *FD) {
return ScannedDecls[FD]; }
2789 bool WasAlreadyScannedInCurrentStage(
const FunctionDecl *FD) {
2790 return WasAlreadyScannedInCurrentStage(GetScannedStages(FD));
2793 bool WasAlreadyScannedInCurrentStage(
unsigned ScannerStages) {
2794 return ScannerStages & CurrentShaderStageBit;
2797 static bool NeverBeenScanned(
unsigned ScannedStages) {
2798 return ScannedStages == 0;
2802 void HandleFunctionOrMethodRef(FunctionDecl *FD, Expr *RefExpr);
2803 void CheckDeclAvailability(NamedDecl *D,
const AvailabilityAttr *AA,
2805 const AvailabilityAttr *FindAvailabilityAttr(
const Decl *D);
2806 bool HasMatchingEnvironmentOrNone(
const AvailabilityAttr *AA);
2809 DiagnoseHLSLAvailability(Sema &SemaRef)
2811 CurrentShaderEnvironment(llvm::Triple::UnknownEnvironment),
2812 CurrentShaderStageBit(0), ReportOnlyShaderStageIssues(
false) {}
2815 void RunOnTranslationUnit(
const TranslationUnitDecl *TU);
2816 void RunOnFunction(
const FunctionDecl *FD);
2818 bool VisitDeclRefExpr(DeclRefExpr *DRE)
override {
2819 FunctionDecl *FD = llvm::dyn_cast<FunctionDecl>(DRE->
getDecl());
2821 HandleFunctionOrMethodRef(FD, DRE);
2825 bool VisitMemberExpr(MemberExpr *ME)
override {
2826 FunctionDecl *FD = llvm::dyn_cast<FunctionDecl>(ME->
getMemberDecl());
2828 HandleFunctionOrMethodRef(FD, ME);
2833void DiagnoseHLSLAvailability::HandleFunctionOrMethodRef(
FunctionDecl *FD,
2836 "expected DeclRefExpr or MemberExpr");
2840 if (FD->
hasBody(FDWithBody)) {
2841 if (!WasAlreadyScannedInCurrentStage(FDWithBody))
2842 DeclsToScan.push_back(FDWithBody);
2847 const AvailabilityAttr *AA = FindAvailabilityAttr(FD);
2849 CheckDeclAvailability(
2853void DiagnoseHLSLAvailability::RunOnTranslationUnit(
2862 DeclContextsToScan.push_back(TU);
2864 while (!DeclContextsToScan.empty()) {
2865 const DeclContext *DC = DeclContextsToScan.pop_back_val();
2866 for (
auto &D : DC->
decls()) {
2873 if (llvm::dyn_cast<NamespaceDecl>(D) || llvm::dyn_cast<ExportDecl>(D)) {
2874 DeclContextsToScan.push_back(llvm::dyn_cast<DeclContext>(D));
2879 const FunctionDecl *FD = llvm::dyn_cast<FunctionDecl>(D);
2884 if (HLSLShaderAttr *ShaderAttr = FD->
getAttr<HLSLShaderAttr>()) {
2885 SetShaderStageContext(ShaderAttr->getType());
2894 for (
const auto *Redecl : FD->
redecls()) {
2895 if (Redecl->isInExportDeclContext()) {
2902 SetUnknownShaderStageContext();
2910void DiagnoseHLSLAvailability::RunOnFunction(
const FunctionDecl *FD) {
2911 assert(DeclsToScan.empty() &&
"DeclsToScan should be empty");
2912 DeclsToScan.push_back(FD);
2914 while (!DeclsToScan.empty()) {
2922 const unsigned ScannedStages = GetScannedStages(FD);
2923 if (WasAlreadyScannedInCurrentStage(ScannedStages))
2926 ReportOnlyShaderStageIssues = !NeverBeenScanned(ScannedStages);
2928 AddToScannedFunctions(FD);
2933bool DiagnoseHLSLAvailability::HasMatchingEnvironmentOrNone(
2934 const AvailabilityAttr *AA) {
2939 llvm::Triple::EnvironmentType CurrentEnv = GetCurrentShaderEnvironment();
2940 if (CurrentEnv == llvm::Triple::UnknownEnvironment)
2943 llvm::Triple::EnvironmentType AttrEnv =
2944 AvailabilityAttr::getEnvironmentType(IIEnvironment->
getName());
2946 return CurrentEnv == AttrEnv;
2949const AvailabilityAttr *
2950DiagnoseHLSLAvailability::FindAvailabilityAttr(
const Decl *D) {
2951 AvailabilityAttr
const *PartialMatch =
nullptr;
2955 for (
const auto *A : D->
attrs()) {
2956 if (
const auto *Avail = dyn_cast<AvailabilityAttr>(A)) {
2957 const AvailabilityAttr *EffectiveAvail = Avail->getEffectiveAttr();
2958 StringRef AttrPlatform = EffectiveAvail->getPlatform()->getName();
2959 StringRef TargetPlatform =
2963 if (AttrPlatform == TargetPlatform) {
2965 if (HasMatchingEnvironmentOrNone(EffectiveAvail))
2967 PartialMatch = Avail;
2971 return PartialMatch;
2976void DiagnoseHLSLAvailability::CheckDeclAvailability(
NamedDecl *D,
2977 const AvailabilityAttr *AA,
2996 if (ReportOnlyShaderStageIssues)
3002 if (InUnknownShaderStageContext())
3007 bool EnvironmentMatches = HasMatchingEnvironmentOrNone(AA);
3008 VersionTuple Introduced = AA->getIntroduced();
3017 llvm::StringRef PlatformName(
3020 llvm::StringRef CurrentEnvStr =
3021 llvm::Triple::getEnvironmentTypeName(GetCurrentShaderEnvironment());
3023 llvm::StringRef AttrEnvStr =
3024 AA->getEnvironment() ? AA->getEnvironment()->getName() :
"";
3025 bool UseEnvironment = !AttrEnvStr.empty();
3027 if (EnvironmentMatches) {
3028 SemaRef.
Diag(
Range.getBegin(), diag::warn_hlsl_availability)
3029 <<
Range << D << PlatformName << Introduced.getAsString()
3030 << UseEnvironment << CurrentEnvStr;
3032 SemaRef.
Diag(
Range.getBegin(), diag::warn_hlsl_availability_unavailable)
3036 SemaRef.
Diag(D->
getLocation(), diag::note_partial_availability_specified_here)
3037 << D << PlatformName << Introduced.getAsString()
3039 << UseEnvironment << AttrEnvStr << CurrentEnvStr;
3046 if (!DefaultCBufferDecls.empty()) {
3049 DefaultCBufferDecls);
3052 SemaRef.getCurLexicalContext()->addDecl(DefaultCBuffer);
3056 for (
const Decl *VD : DefaultCBufferDecls) {
3057 const HLSLResourceBindingAttr *RBA =
3058 VD->
getAttr<HLSLResourceBindingAttr>();
3059 if (RBA && RBA->hasRegisterSlot() &&
3060 RBA->getRegisterType() == HLSLResourceBindingAttr::RegisterType::C) {
3067 SemaRef.Consumer.HandleTopLevelDecl(DG);
3069 diagnoseAvailabilityViolations(TU);
3078 "expected member expr to have resource record type or array of them");
3084 const Expr *NonConstIndexExpr =
nullptr;
3087 if (
const DeclRefExpr *DRE = dyn_cast<DeclRefExpr>(E)) {
3088 if (!NonConstIndexExpr)
3096 diag::err_hlsl_resource_member_array_access_not_constant);
3100 if (
const auto *ASE = dyn_cast<ArraySubscriptExpr>(E)) {
3101 const Expr *IdxExpr = ASE->getIdx();
3103 NonConstIndexExpr = IdxExpr;
3105 }
else if (
const auto *SubME = dyn_cast<MemberExpr>(E)) {
3106 E = SubME->getBase();
3107 }
else if (
const auto *ICE = dyn_cast<ImplicitCastExpr>(E)) {
3108 E = ICE->getSubExpr();
3110 llvm_unreachable(
"unexpected expr type in resource member access");
3123 TI.
getTriple().getEnvironment() != llvm::Triple::EnvironmentType::Library)
3126 DiagnoseHLSLAvailability(
SemaRef).RunOnTranslationUnit(TU);
3133 for (
unsigned I = 1, N = TheCall->
getNumArgs(); I < N; ++I) {
3136 S->
Diag(TheCall->
getBeginLoc(), diag::err_vec_builtin_incompatible_vector)
3161 for (
unsigned I = 0; I < TheCall->
getNumArgs(); ++I) {
3176 if (!BaseType->isFloat32Type())
3177 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3178 << ArgOrdinal << 5 << 0
3190 if (!BaseType->isHalfType() && !BaseType->isFloat32Type())
3191 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3192 << ArgOrdinal << 5 << 0
3206 if (!BaseType->isDoubleType()) {
3209 return S->
Diag(Loc, diag::err_builtin_requires_double_type)
3210 << ArgOrdinal << PassedType;
3217 unsigned ArgIndex) {
3218 auto *Arg = TheCall->
getArg(ArgIndex);
3220 if (Arg->IgnoreCasts()->isModifiableLvalue(S->
Context, &OrigLoc) ==
3223 S->
Diag(OrigLoc, diag::error_hlsl_inout_lvalue) << Arg << 0;
3233 if (VecTy->getElementType()->isDoubleType())
3234 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3235 << ArgOrdinal << 1 << 0 << 1
3245 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3246 << ArgOrdinal << 5 << 1
3255 if (VecTy->getElementType()->isUnsignedIntegerType())
3258 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3259 << ArgOrdinal << 4 << 3 << 0
3268 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3269 << ArgOrdinal << 5 << 3
3275 unsigned ArgOrdinal,
unsigned Width) {
3278 ArgTy = VTy->getElementType();
3280 uint64_t ElementBitCount =
3282 if (ElementBitCount != Width) {
3284 diag::err_integer_incorrect_bit_count)
3285 << Width << ElementBitCount;
3302 unsigned ArgIndex) {
3311 diag::err_typecheck_expect_scalar_or_vector)
3312 << ArgType << Scalar;
3319 QualType Scalar,
unsigned ArgIndex) {
3330 if (
const auto *VTy = ArgType->getAs<
VectorType>()) {
3343 diag::err_typecheck_expect_scalar_or_vector_or_matrix)
3344 << ArgType << Scalar;
3349 unsigned ArgIndex) {
3354 if (!(ArgType->isScalarType() ||
3355 (VTy && VTy->getElementType()->isScalarType()))) {
3357 diag::err_typecheck_expect_any_scalar_or_vector)
3367 unsigned ArgIndex) {
3369 assert(ArgIndex < TheCall->getNumArgs());
3377 diag::err_typecheck_expect_any_scalar_or_vector)
3402 diag::err_typecheck_call_different_arg_types)
3421 Arg1ScalarTy = VTy->getElementType();
3425 Arg2ScalarTy = VTy->getElementType();
3428 S->
Diag(Arg1->
getBeginLoc(), diag::err_hlsl_builtin_scalar_vector_mismatch)
3429 << 1 << TheCall->
getCallee() << Arg1Ty << Arg2Ty;
3439 if (Arg1Length > 0 && Arg0Length != Arg1Length) {
3441 diag::err_typecheck_vector_lengths_not_equal)
3447 if (Arg2Length > 0 && Arg0Length != Arg2Length) {
3449 diag::err_typecheck_vector_lengths_not_equal)
3461 assert(TheCall->
getNumArgs() > IndexArgIndex &&
"Index argument missing");
3464 unsigned int ActualDim = 1;
3466 ActualDim = VTy->getNumElements();
3467 IndexTy = VTy->getElementType();
3471 diag::err_typecheck_expect_int)
3477 const HLSLAttributedResourceType *ResTy =
3479 assert(ResTy &&
"Resource argument must be a resource");
3480 HLSLAttributedResourceType::Attributes ResAttrs = ResTy->getAttrs();
3482 unsigned int ExpectedDim = 1;
3483 if (ResAttrs.ResourceDimension != llvm::dxil::ResourceDimension::Unknown)
3486 if (ActualDim != ExpectedDim) {
3488 diag::err_hlsl_builtin_resource_coordinate_dimension_mismatch)
3499 llvm::function_ref<
bool(
const HLSLAttributedResourceType *ResType)> Check =
3503 const HLSLAttributedResourceType *ResTy =
3507 diag::err_typecheck_expect_hlsl_resource)
3511 if (Check && Check(ResTy)) {
3513 diag::err_invalid_hlsl_resource_type)
3521 QualType BaseType,
unsigned ExpectedCount,
3523 unsigned PassedCount = 1;
3525 PassedCount = VecTy->getNumElements();
3527 if (PassedCount != ExpectedCount) {
3530 S->
Diag(Loc, diag::err_typecheck_convert_incompatible)
3542 [](
const HLSLAttributedResourceType *ResType) {
3543 return ResType->getAttrs().ResourceDimension ==
3544 llvm::dxil::ResourceDimension::Unknown;
3550 [](
const HLSLAttributedResourceType *ResType) {
3551 return ResType->getAttrs().ResourceClass !=
3552 llvm::hlsl::ResourceClass::Sampler;
3560 unsigned ExpectedDim =
3588 unsigned NextIdx = 3;
3594 diag::err_typecheck_convert_incompatible)
3602 Expr *ComponentArg = TheCall->
getArg(NextIdx);
3606 diag::err_typecheck_convert_incompatible)
3613 std::optional<llvm::APSInt> ComponentOpt =
3616 int64_t ComponentVal = ComponentOpt->getSExtValue();
3617 if (ComponentVal != 0) {
3620 assert(ComponentVal >= 0 && ComponentVal <= 3 &&
3621 "The component is not in the expected range.");
3623 diag::err_hlsl_gathercmp_invalid_component)
3633 const HLSLAttributedResourceType *ResourceTy =
3636 unsigned ExpectedDim =
3645 assert(ResourceTy->hasContainedType() &&
3646 "Expecting a contained type for resource with a dimension "
3648 QualType ReturnType = ResourceTy->getContainedType();
3652 S.
Diag(TheCall->
getBeginLoc(), diag::err_hlsl_samplecmp_requires_float);
3658 ReturnType = VecTy->getElementType();
3671 [](
const HLSLAttributedResourceType *ResType) {
3672 return ResType->getAttrs().ResourceDimension ==
3673 llvm::dxil::ResourceDimension::Unknown;
3681 unsigned ExpectedDim =
3690 EltTy = VTy->getElementType();
3705 TheCall->
setType(ResourceTy->getContainedType());
3710 unsigned MinArgs, MaxArgs;
3738 const HLSLAttributedResourceType *ResourceTy =
3740 unsigned ExpectedDim =
3743 unsigned NextIdx = 3;
3752 diag::err_typecheck_convert_incompatible)
3787 diag::err_typecheck_convert_incompatible)
3793 assert(ResourceTy->hasContainedType() &&
3794 "Expecting a contained type for resource with a dimension "
3796 QualType ReturnType = ResourceTy->getContainedType();
3799 S.
Diag(TheCall->
getBeginLoc(), diag::err_hlsl_samplecmp_requires_float);
3812 switch (BuiltinID) {
3813 case Builtin::BI__builtin_hlsl_adduint64: {
3814 if (
SemaRef.checkArgCount(TheCall, 2))
3828 if (NumElementsArg != 2 && NumElementsArg != 4) {
3830 << 1 << 64 << NumElementsArg * 32;
3844 case Builtin::BI__builtin_hlsl_resource_getpointer: {
3845 if (
SemaRef.checkArgCount(TheCall, 2) ||
3852 QualType ContainedTy = ResourceTy->getContainedType();
3855 ReturnType =
SemaRef.Context.getPointerType(ReturnType);
3861 case Builtin::BI__builtin_hlsl_resource_getpointer_typed: {
3862 if (
SemaRef.checkArgCount(TheCall, 3) ||
3869 "expected pointer type for second argument");
3876 diag::err_invalid_use_of_array_type);
3880 ReturnType =
SemaRef.Context.getPointerType(ReturnType);
3885 case Builtin::BI__builtin_hlsl_resource_load_with_status: {
3886 if (
SemaRef.checkArgCount(TheCall, 3) ||
3889 SemaRef.getASTContext().UnsignedIntTy) ||
3891 SemaRef.getASTContext().UnsignedIntTy) ||
3897 QualType ReturnType = ResourceTy->getContainedType();
3902 case Builtin::BI__builtin_hlsl_resource_load_with_status_typed: {
3903 if (
SemaRef.checkArgCount(TheCall, 4) ||
3906 SemaRef.getASTContext().UnsignedIntTy) ||
3908 SemaRef.getASTContext().UnsignedIntTy) ||
3914 "expected pointer type for second argument");
3921 diag::err_invalid_use_of_array_type);
3927 case Builtin::BI__builtin_hlsl_resource_load_level:
3929 case Builtin::BI__builtin_hlsl_resource_sample:
3931 case Builtin::BI__builtin_hlsl_resource_sample_bias:
3933 case Builtin::BI__builtin_hlsl_resource_sample_grad:
3935 case Builtin::BI__builtin_hlsl_resource_sample_level:
3937 case Builtin::BI__builtin_hlsl_resource_sample_cmp:
3939 case Builtin::BI__builtin_hlsl_resource_sample_cmp_level_zero:
3941 case Builtin::BI__builtin_hlsl_resource_calculate_lod:
3942 case Builtin::BI__builtin_hlsl_resource_calculate_lod_unclamped:
3944 case Builtin::BI__builtin_hlsl_resource_gather:
3946 case Builtin::BI__builtin_hlsl_resource_gather_cmp:
3948 case Builtin::BI__builtin_hlsl_resource_uninitializedhandle: {
3949 assert(TheCall->
getNumArgs() == 1 &&
"expected 1 arg");
3955 case Builtin::BI__builtin_hlsl_resource_handlefrombinding: {
3956 assert(TheCall->
getNumArgs() == 6 &&
"expected 6 args");
3962 case Builtin::BI__builtin_hlsl_resource_handlefromimplicitbinding: {
3963 assert(TheCall->
getNumArgs() == 6 &&
"expected 6 args");
3969 case Builtin::BI__builtin_hlsl_resource_counterhandlefromimplicitbinding: {
3970 assert(TheCall->
getNumArgs() == 3 &&
"expected 3 args");
3973 auto *MainResType = MainHandleTy->
getAs<HLSLAttributedResourceType>();
3974 auto MainAttrs = MainResType->getAttrs();
3975 assert(!MainAttrs.IsCounter &&
"cannot create a counter from a counter");
3976 MainAttrs.IsCounter =
true;
3978 MainResType->getWrappedType(), MainResType->getContainedType(),
3982 TheCall->
setType(CounterHandleTy);
3985 case Builtin::BI__builtin_hlsl_and:
3986 case Builtin::BI__builtin_hlsl_or: {
3987 if (
SemaRef.checkArgCount(TheCall, 2))
4001 case Builtin::BI__builtin_hlsl_all:
4002 case Builtin::BI__builtin_hlsl_any: {
4003 if (
SemaRef.checkArgCount(TheCall, 1))
4009 case Builtin::BI__builtin_hlsl_asdouble: {
4010 if (
SemaRef.checkArgCount(TheCall, 2))
4014 SemaRef.Context.UnsignedIntTy,
4019 SemaRef.Context.UnsignedIntTy,
4028 case Builtin::BI__builtin_hlsl_elementwise_clamp: {
4029 if (
SemaRef.BuiltinElementwiseTernaryMath(
4035 case Builtin::BI__builtin_hlsl_dot: {
4037 if (
SemaRef.BuiltinVectorToScalarMath(TheCall))
4043 case Builtin::BI__builtin_hlsl_elementwise_firstbithigh:
4044 case Builtin::BI__builtin_hlsl_elementwise_firstbitlow: {
4045 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
4055 EltTy = VecTy->getElementType();
4056 ResTy =
SemaRef.Context.getExtVectorType(ResTy, VecTy->getNumElements());
4069 case Builtin::BI__builtin_hlsl_select: {
4070 if (
SemaRef.checkArgCount(TheCall, 3))
4078 if (VTy && VTy->getElementType()->isBooleanType() &&
4083 case Builtin::BI__builtin_hlsl_elementwise_saturate:
4084 case Builtin::BI__builtin_hlsl_elementwise_rcp: {
4085 if (
SemaRef.checkArgCount(TheCall, 1))
4091 diag::err_builtin_invalid_arg_type)
4094 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
4098 case Builtin::BI__builtin_hlsl_elementwise_degrees:
4099 case Builtin::BI__builtin_hlsl_elementwise_radians:
4100 case Builtin::BI__builtin_hlsl_elementwise_rsqrt:
4101 case Builtin::BI__builtin_hlsl_elementwise_frac:
4102 case Builtin::BI__builtin_hlsl_elementwise_ddx_coarse:
4103 case Builtin::BI__builtin_hlsl_elementwise_ddy_coarse:
4104 case Builtin::BI__builtin_hlsl_elementwise_ddx_fine:
4105 case Builtin::BI__builtin_hlsl_elementwise_ddy_fine: {
4106 if (
SemaRef.checkArgCount(TheCall, 1))
4111 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
4115 case Builtin::BI__builtin_hlsl_elementwise_isinf:
4116 case Builtin::BI__builtin_hlsl_elementwise_isnan: {
4117 if (
SemaRef.checkArgCount(TheCall, 1))
4122 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
4127 case Builtin::BI__builtin_hlsl_lerp: {
4128 if (
SemaRef.checkArgCount(TheCall, 3))
4135 if (
SemaRef.BuiltinElementwiseTernaryMath(TheCall))
4139 case Builtin::BI__builtin_hlsl_mad: {
4140 if (
SemaRef.BuiltinElementwiseTernaryMath(
4146 case Builtin::BI__builtin_hlsl_mul: {
4147 if (
SemaRef.checkArgCount(TheCall, 2))
4156 if (
const auto *VTy = T->getAs<
VectorType>())
4157 return VTy->getElementType();
4159 return MTy->getElementType();
4163 QualType EltTy0 = getElemType(Ty0);
4172 if (IsVec0 && IsMat1) {
4175 }
else if (IsMat0 && IsVec1) {
4179 assert(IsMat0 && IsMat1);
4189 case Builtin::BI__builtin_hlsl_normalize: {
4190 if (
SemaRef.checkArgCount(TheCall, 1))
4201 case Builtin::BI__builtin_elementwise_fma: {
4202 if (
SemaRef.checkArgCount(TheCall, 3) ||
4217 case Builtin::BI__builtin_hlsl_transpose: {
4218 if (
SemaRef.checkArgCount(TheCall, 1))
4227 << 1 << 3 << 0 << 0 << ArgTy;
4232 MatTy->getElementType(), MatTy->getNumColumns(), MatTy->getNumRows());
4236 case Builtin::BI__builtin_hlsl_elementwise_sign: {
4237 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
4245 case Builtin::BI__builtin_hlsl_step: {
4246 if (
SemaRef.checkArgCount(TheCall, 2))
4258 case Builtin::BI__builtin_hlsl_wave_active_all_equal: {
4259 if (
SemaRef.checkArgCount(TheCall, 1))
4273 unsigned NumElts = VecTy->getNumElements();
4283 case Builtin::BI__builtin_hlsl_wave_active_max:
4284 case Builtin::BI__builtin_hlsl_wave_active_min:
4285 case Builtin::BI__builtin_hlsl_wave_active_sum:
4286 case Builtin::BI__builtin_hlsl_wave_active_product: {
4287 if (
SemaRef.checkArgCount(TheCall, 1))
4300 case Builtin::BI__builtin_hlsl_wave_active_bit_or:
4301 case Builtin::BI__builtin_hlsl_wave_active_bit_xor:
4302 case Builtin::BI__builtin_hlsl_wave_active_bit_and: {
4303 if (
SemaRef.checkArgCount(TheCall, 1))
4318 (VTy && VTy->getElementType()->isIntegerType()))) {
4320 diag::err_builtin_invalid_arg_type)
4321 << ArgTyExpr <<
SemaRef.Context.UnsignedIntTy << 1 << 0 << 0;
4331 case Builtin::BI__builtin_elementwise_bitreverse: {
4339 case Builtin::BI__builtin_hlsl_wave_prefix_count_bits: {
4340 if (
SemaRef.checkArgCount(TheCall, 1))
4345 if (!(
ArgType->isScalarType())) {
4347 diag::err_typecheck_expect_any_scalar_or_vector)
4352 if (!(
ArgType->isBooleanType())) {
4354 diag::err_typecheck_expect_any_scalar_or_vector)
4361 case Builtin::BI__builtin_hlsl_wave_read_lane_at: {
4362 if (
SemaRef.checkArgCount(TheCall, 2))
4370 diag::err_typecheck_convert_incompatible)
4371 << ArgTyIndex <<
SemaRef.Context.UnsignedIntTy << 1 << 0 << 0;
4384 case Builtin::BI__builtin_hlsl_wave_get_lane_index: {
4385 if (
SemaRef.checkArgCount(TheCall, 0))
4389 case Builtin::BI__builtin_hlsl_wave_prefix_sum:
4390 case Builtin::BI__builtin_hlsl_wave_prefix_product: {
4391 if (
SemaRef.checkArgCount(TheCall, 1))
4404 case Builtin::BI__builtin_hlsl_quad_read_across_x:
4405 case Builtin::BI__builtin_hlsl_quad_read_across_y: {
4406 if (
SemaRef.checkArgCount(TheCall, 1))
4418 case Builtin::BI__builtin_hlsl_elementwise_splitdouble: {
4419 if (
SemaRef.checkArgCount(TheCall, 3))
4434 case Builtin::BI__builtin_hlsl_elementwise_clip: {
4435 if (
SemaRef.checkArgCount(TheCall, 1))
4442 case Builtin::BI__builtin_elementwise_acos:
4443 case Builtin::BI__builtin_elementwise_asin:
4444 case Builtin::BI__builtin_elementwise_atan:
4445 case Builtin::BI__builtin_elementwise_atan2:
4446 case Builtin::BI__builtin_elementwise_ceil:
4447 case Builtin::BI__builtin_elementwise_cos:
4448 case Builtin::BI__builtin_elementwise_cosh:
4449 case Builtin::BI__builtin_elementwise_exp:
4450 case Builtin::BI__builtin_elementwise_exp2:
4451 case Builtin::BI__builtin_elementwise_exp10:
4452 case Builtin::BI__builtin_elementwise_floor:
4453 case Builtin::BI__builtin_elementwise_fmod:
4454 case Builtin::BI__builtin_elementwise_log:
4455 case Builtin::BI__builtin_elementwise_log2:
4456 case Builtin::BI__builtin_elementwise_log10:
4457 case Builtin::BI__builtin_elementwise_pow:
4458 case Builtin::BI__builtin_elementwise_roundeven:
4459 case Builtin::BI__builtin_elementwise_sin:
4460 case Builtin::BI__builtin_elementwise_sinh:
4461 case Builtin::BI__builtin_elementwise_sqrt:
4462 case Builtin::BI__builtin_elementwise_tan:
4463 case Builtin::BI__builtin_elementwise_tanh:
4464 case Builtin::BI__builtin_elementwise_trunc: {
4470 case Builtin::BI__builtin_hlsl_buffer_update_counter: {
4471 assert(TheCall->
getNumArgs() == 2 &&
"expected 2 args");
4472 auto checkResTy = [](
const HLSLAttributedResourceType *ResTy) ->
bool {
4473 return !(ResTy->getAttrs().ResourceClass == ResourceClass::UAV &&
4474 ResTy->getAttrs().RawBuffer && ResTy->hasContainedType());
4479 std::optional<llvm::APSInt> Offset =
4481 if (!Offset.has_value() ||
std::abs(Offset->getExtValue()) != 1) {
4483 diag::err_hlsl_expect_arg_const_int_one_or_neg_one)
4489 case Builtin::BI__builtin_hlsl_elementwise_f16tof32: {
4490 if (
SemaRef.checkArgCount(TheCall, 1))
4501 ArgTy = VTy->getElementType();
4504 diag::err_builtin_invalid_arg_type)
4513 case Builtin::BI__builtin_hlsl_elementwise_f32tof16: {
4514 if (
SemaRef.checkArgCount(TheCall, 1))
4529 WorkList.push_back(BaseTy);
4530 while (!WorkList.empty()) {
4531 QualType T = WorkList.pop_back_val();
4532 T = T.getCanonicalType().getUnqualifiedType();
4533 if (
const auto *AT = dyn_cast<ConstantArrayType>(T)) {
4541 for (uint64_t Ct = 0; Ct < AT->
getZExtSize(); ++Ct)
4542 llvm::append_range(List, ElementFields);
4547 if (
const auto *VT = dyn_cast<VectorType>(T)) {
4548 List.insert(List.end(), VT->getNumElements(), VT->getElementType());
4551 if (
const auto *MT = dyn_cast<ConstantMatrixType>(T)) {
4552 List.insert(List.end(), MT->getNumElementsFlattened(),
4553 MT->getElementType());
4556 if (
const auto *RD = T->getAsCXXRecordDecl()) {
4557 if (RD->isStandardLayout())
4558 RD = RD->getStandardLayoutBaseWithFields();
4562 if (RD->
isUnion() || !RD->isAggregate()) {
4568 for (
const auto *FD : RD->
fields())
4569 if (!FD->isUnnamedBitField())
4570 FieldTypes.push_back(FD->
getType());
4572 std::reverse(FieldTypes.begin(), FieldTypes.end());
4573 llvm::append_range(WorkList, FieldTypes);
4577 if (!RD->isStandardLayout()) {
4579 for (
const auto &
Base : RD->bases())
4580 FieldTypes.push_back(
Base.getType());
4581 std::reverse(FieldTypes.begin(), FieldTypes.end());
4582 llvm::append_range(WorkList, FieldTypes);
4604 if (
SemaRef.Context.getTypeSize(QT) / 8 > 16)
4610 int ArraySize = VT->getNumElements();
4615 QualType ElTy = VT->getElementType();
4619 if (
SemaRef.Context.getTypeSize(QT) / 8 > 16)
4635 if (
SemaRef.getASTContext().hasSameType(T1, T2))
4644 return llvm::equal(T1Types, T2Types,
4646 return SemaRef.IsLayoutCompatible(LHS, RHS);
4655 bool HadError =
false;
4657 for (
unsigned i = 0, e =
New->getNumParams(); i != e; ++i) {
4665 const auto *NDAttr = NewParam->
getAttr<HLSLParamModifierAttr>();
4666 unsigned NSpellingIdx = (NDAttr ? NDAttr->getSpellingListIndex() : 0);
4667 const auto *ODAttr = OldParam->
getAttr<HLSLParamModifierAttr>();
4668 unsigned OSpellingIdx = (ODAttr ? ODAttr->getSpellingListIndex() : 0);
4670 if (NSpellingIdx != OSpellingIdx) {
4672 diag::err_hlsl_param_qualifier_mismatch)
4673 << NDAttr << NewParam;
4689 if (
SemaRef.getASTContext().hasSameUnqualifiedType(SrcTy, DestTy))
4704 llvm_unreachable(
"HLSL doesn't support pointers.");
4707 llvm_unreachable(
"HLSL doesn't support complex types.");
4709 llvm_unreachable(
"HLSL doesn't support fixed point types.");
4711 llvm_unreachable(
"Should have returned before this");
4721 llvm_unreachable(
"HLSL doesn't support complex types.");
4723 llvm_unreachable(
"HLSL doesn't support fixed point types.");
4728 llvm_unreachable(
"HLSL doesn't support pointers.");
4730 llvm_unreachable(
"Should have returned before this");
4736 llvm_unreachable(
"HLSL doesn't support pointers.");
4739 llvm_unreachable(
"HLSL doesn't support fixed point types.");
4743 llvm_unreachable(
"HLSL doesn't support complex types.");
4746 llvm_unreachable(
"Unhandled scalar cast");
4767 !(SrcMatTy && SrcMatTy->getNumElementsFlattened() == 1))
4773 SrcTy = SrcMatTy->getElementType();
4778 for (
unsigned I = 0, Size = DestTypes.size(); I < Size; ++I) {
4779 if (DestTypes[I]->isUnionType())
4811 if (SrcTypes.size() < DestTypes.size())
4814 unsigned SrcSize = SrcTypes.size();
4815 unsigned DstSize = DestTypes.size();
4817 for (I = 0; I < DstSize && I < SrcSize; I++) {
4818 if (SrcTypes[I]->isUnionType() || DestTypes[I]->isUnionType())
4826 for (; I < SrcSize; I++) {
4827 if (SrcTypes[I]->isUnionType())
4834 assert(Param->hasAttr<HLSLParamModifierAttr>() &&
4835 "We should not get here without a parameter modifier expression");
4836 const auto *
Attr = Param->getAttr<HLSLParamModifierAttr>();
4843 << Arg << (IsInOut ? 1 : 0);
4849 QualType Ty = Param->getType().getNonLValueExprType(Ctx);
4856 << Arg << (IsInOut ? 1 : 0);
4868 SemaRef.PerformCopyInitialization(Entity, Param->getBeginLoc(), ArgOpV);
4874 auto *OpV =
new (Ctx)
4879 Res =
SemaRef.ActOnBinOp(
SemaRef.getCurScope(), Param->getBeginLoc(),
4880 tok::equal, ArgOpV, OpV);
4896 "Pointer and reference types cannot be inout or out parameters");
4897 Ty =
SemaRef.getASTContext().getLValueReferenceType(Ty);
4913 for (
const auto *FD : RD->
fields()) {
4917 assert(RD->getNumBases() <= 1 &&
4918 "HLSL doesn't support multiple inheritance");
4919 return RD->getNumBases()
4924 if (
const auto *AT = dyn_cast<ArrayType>(Ty)) {
4925 if (
const auto *CAT = dyn_cast<ConstantArrayType>(AT))
4937 bool IsVKPushConstant = IsVulkan && VD->
hasAttr<HLSLVkPushConstantAttr>();
4942 !VD->
hasAttr<HLSLVkConstantIdAttr>() && !IsVKPushConstant &&
4948 if (
Decl->getType().hasAddressSpace())
4951 if (
Decl->getType()->isDependentType())
4963 if (
Decl->
hasAttr<HLSLVkExtBuiltinOutputAttr>()) {
4977 llvm::Triple::Vulkan;
4978 if (IsVulkan &&
Decl->
hasAttr<HLSLVkPushConstantAttr>()) {
4979 if (HasDeclaredAPushConstant)
4985 HasDeclaredAPushConstant =
true;
5012class StructBindingContext {
5015 HLSLResourceBindingAttr *RegBindingsAttrs[4];
5016 unsigned RegBindingOffset[4];
5019 static_assert(
static_cast<unsigned>(RegisterType::SRV) == 0 &&
5020 static_cast<unsigned>(RegisterType::UAV) == 1 &&
5021 static_cast<unsigned>(RegisterType::CBuffer) == 2 &&
5022 static_cast<unsigned>(RegisterType::Sampler) == 3,
5023 "unexpected register type values");
5026 HLSLVkBindingAttr *VkBindingAttr;
5027 unsigned VkBindingOffset;
5032 StructBindingContext(
VarDecl *VD) {
5033 for (
unsigned i = 0; i < 4; ++i) {
5034 RegBindingsAttrs[i] =
nullptr;
5035 RegBindingOffset[i] = 0;
5037 VkBindingAttr =
nullptr;
5038 VkBindingOffset = 0;
5044 if (
auto *RBA = dyn_cast<HLSLResourceBindingAttr>(A)) {
5046 unsigned RegTypeIdx =
static_cast<unsigned>(RegType);
5049 RegBindingsAttrs[RegTypeIdx] = RBA;
5054 if (
auto *VBA = dyn_cast<HLSLVkBindingAttr>(A))
5055 VkBindingAttr = VBA;
5062 Attr *createBindingAttr(SemaHLSL &S, ASTContext &AST,
RegisterType RegType,
5064 assert(
static_cast<unsigned>(RegType) < 4 &&
"unexpected register type");
5066 if (VkBindingAttr) {
5067 unsigned Offset = VkBindingOffset;
5068 VkBindingOffset +=
Range;
5069 return HLSLVkBindingAttr::CreateImplicit(
5070 AST, VkBindingAttr->getBinding() + Offset, VkBindingAttr->getSet(),
5071 VkBindingAttr->getRange());
5074 HLSLResourceBindingAttr *RBA =
5075 RegBindingsAttrs[
static_cast<unsigned>(RegType)];
5076 HLSLResourceBindingAttr *NewAttr =
nullptr;
5078 if (RBA && RBA->hasRegisterSlot()) {
5081 unsigned Offset = RegBindingOffset[
static_cast<unsigned>(RegType)];
5082 RegBindingOffset[
static_cast<unsigned>(RegType)] += Range;
5084 unsigned NewSlotNumber = RBA->getSlotNumber() + Offset;
5085 StringRef NewSlotNumberStr =
5087 NewAttr = HLSLResourceBindingAttr::CreateImplicit(
5088 AST, NewSlotNumberStr, RBA->getSpace(), RBA->getRange());
5089 NewAttr->setBinding(RegType, NewSlotNumber, RBA->getSpaceNumber());
5093 NewAttr = HLSLResourceBindingAttr::CreateImplicit(AST,
"",
"0", {});
5094 NewAttr->setBinding(RegType, std::nullopt,
5095 RBA ? RBA->getSpaceNumber() : 0);
5105static void createGlobalResourceDeclForStruct(
5107 QualType ResTy, StructBindingContext &BindingCtx) {
5109 "expected resource type or array of resources");
5119 const HLSLAttributedResourceType *ResHandleTy =
nullptr;
5120 if (
const auto *AT = dyn_cast<ArrayType>(ResTy.
getTypePtr())) {
5121 const auto *CAT = dyn_cast<ConstantArrayType>(AT);
5125 ResHandleTy = HLSLAttributedResourceType::findHandleTypeOnResource(
5129 Attr *BindingAttr = BindingCtx.createBindingAttr(
5131 ResDecl->
addAttr(BindingAttr);
5132 ResDecl->
addAttr(InternalLinkageAttr::CreateImplicit(AST));
5141 HLSLAssociatedResourceDeclAttr::CreateImplicit(AST, ResDecl));
5148static void handleArrayOfStructWithResources(
5150 EmbeddedResourceNameBuilder &NameBuilder, StructBindingContext &BindingCtx);
5155static void handleStructWithResources(
Sema &S,
VarDecl *ParentVD,
5157 EmbeddedResourceNameBuilder &NameBuilder,
5158 StructBindingContext &BindingCtx) {
5161 assert(RD->
getNumBases() <= 1 &&
"HLSL doesn't support multiple inheritance");
5168 handleStructWithResources(S, ParentVD, BaseRD, NameBuilder, BindingCtx);
5182 createGlobalResourceDeclForStruct(S, ParentVD, FD->
getLocation(), II,
5185 handleStructWithResources(S, ParentVD, RD, NameBuilder, BindingCtx);
5187 }
else if (
const auto *ArrayTy = dyn_cast<ConstantArrayType>(FDTy)) {
5189 "resource arrays should have been already handled");
5190 handleArrayOfStructWithResources(S, ParentVD, ArrayTy, NameBuilder,
5199handleArrayOfStructWithResources(
Sema &S,
VarDecl *ParentVD,
5201 EmbeddedResourceNameBuilder &NameBuilder,
5202 StructBindingContext &BindingCtx) {
5210 if (!SubCAT && !ElementRD)
5213 for (
unsigned I = 0, E = CAT->
getSize().getZExtValue(); I < E; ++I) {
5216 handleStructWithResources(S, ParentVD, ElementRD, NameBuilder,
5219 handleArrayOfStructWithResources(S, ParentVD, SubCAT, NameBuilder,
5232void SemaHLSL::handleGlobalStructOrArrayOfWithResources(
VarDecl *VD) {
5233 EmbeddedResourceNameBuilder NameBuilder(VD->
getName());
5234 StructBindingContext BindingCtx(VD);
5238 "Expected non-resource struct or array type");
5241 handleStructWithResources(
SemaRef, VD, RD, NameBuilder, BindingCtx);
5245 if (
const auto *CAT = dyn_cast<ConstantArrayType>(VDTy)) {
5246 handleArrayOfStructWithResources(
SemaRef, VD, CAT, NameBuilder, BindingCtx);
5254 if (
SemaRef.RequireCompleteType(
5257 diag::err_typecheck_decl_incomplete_type)) {
5271 DefaultCBufferDecls.push_back(VD);
5276 collectResourceBindingsOnVarDecl(VD);
5278 if (VD->
hasAttr<HLSLVkConstantIdAttr>())
5290 processExplicitBindingsOnDecl(VD);
5328 handleGlobalStructOrArrayOfWithResources(VD);
5332 if (VD->
hasAttr<HLSLGroupSharedAddressSpaceAttr>())
5341 "expected resource record type");
5357 const char *CreateMethodName;
5359 CreateMethodName = HasCounter ?
"__createFromBindingWithImplicitCounter"
5360 :
"__createFromBinding";
5362 CreateMethodName = HasCounter
5363 ?
"__createFromImplicitBindingWithImplicitCounter"
5364 :
"__createFromImplicitBinding";
5379 Args.push_back(RegSlot);
5387 Args.push_back(OrderId);
5393 Args.push_back(Space);
5397 Args.push_back(RangeSize);
5401 Args.push_back(Index);
5403 StringRef VarName = VD->
getName();
5411 Args.push_back(NameCast);
5419 Args.push_back(CounterId);
5442 SemaRef.CheckCompleteVariableDeclaration(VD);
5448 "expected array of resource records");
5469 lookupMethod(
SemaRef, ResourceDecl,
5470 HasCounter ?
"__createFromBindingWithImplicitCounter"
5471 :
"__createFromBinding",
5475 CreateMethod = lookupMethod(
5477 HasCounter ?
"__createFromImplicitBindingWithImplicitCounter"
5478 :
"__createFromImplicitBinding",
5510std::optional<const DeclBindingInfo *> SemaHLSL::inferGlobalBinding(
Expr *E) {
5511 if (
auto *Ternary = dyn_cast<ConditionalOperator>(E)) {
5512 auto TrueInfo = inferGlobalBinding(Ternary->getTrueExpr());
5513 auto FalseInfo = inferGlobalBinding(Ternary->getFalseExpr());
5514 if (!TrueInfo || !FalseInfo)
5515 return std::nullopt;
5516 if (*TrueInfo != *FalseInfo)
5517 return std::nullopt;
5521 if (
auto *ASE = dyn_cast<ArraySubscriptExpr>(E))
5530 if (
const auto *AttrResType =
5531 HLSLAttributedResourceType::findHandleTypeOnResource(Ty)) {
5533 return Bindings.getDeclBindingInfo(VD, RC);
5540void SemaHLSL::trackLocalResource(
VarDecl *VD,
Expr *E) {
5541 std::optional<const DeclBindingInfo *> ExprBinding = inferGlobalBinding(E);
5544 diag::warn_hlsl_assigning_local_resource_is_not_unique)
5549 if (*ExprBinding ==
nullptr)
5552 auto PrevBinding = Assigns.find(VD);
5553 if (PrevBinding == Assigns.end()) {
5555 Assigns.insert({VD, *ExprBinding});
5560 if (*ExprBinding != PrevBinding->second) {
5562 diag::warn_hlsl_assigning_local_resource_is_not_unique)
5564 SemaRef.Diag(VD->getLocation(), diag::note_var_declared_here) << VD;
5575 "expected LHS to be a resource record or array of resource records");
5576 if (Opc != BO_Assign)
5581 while (
auto *ASE = dyn_cast<ArraySubscriptExpr>(E))
5589 SemaRef.Diag(Loc, diag::err_hlsl_assign_to_global_resource) << VD;
5594 trackLocalResource(VD, RHSExpr);
5602void SemaHLSL::collectResourceBindingsOnVarDecl(
VarDecl *VD) {
5604 "expected global variable that contains HLSL resource");
5607 if (
const HLSLBufferDecl *CBufferOrTBuffer = dyn_cast<HLSLBufferDecl>(VD)) {
5608 Bindings.addDeclBindingInfo(VD, CBufferOrTBuffer->isCBuffer()
5609 ? ResourceClass::CBuffer
5610 : ResourceClass::SRV);
5623 if (
const HLSLAttributedResourceType *AttrResType =
5624 HLSLAttributedResourceType::findHandleTypeOnResource(Ty)) {
5625 Bindings.addDeclBindingInfo(VD, AttrResType->getAttrs().ResourceClass);
5630 if (
const RecordType *RT = dyn_cast<RecordType>(Ty))
5631 collectResourceBindingsOnUserRecordDecl(VD, RT);
5637void SemaHLSL::processExplicitBindingsOnDecl(
VarDecl *VD) {
5640 bool HasBinding =
false;
5641 for (Attr *A : VD->
attrs()) {
5644 if (
auto PA = VD->
getAttr<HLSLVkPushConstantAttr>())
5645 Diag(PA->getLoc(), diag::err_hlsl_attr_incompatible) << A << PA;
5648 HLSLResourceBindingAttr *RBA = dyn_cast<HLSLResourceBindingAttr>(A);
5649 if (!RBA || !RBA->hasRegisterSlot())
5654 assert(RT != RegisterType::I &&
"invalid or obsolete register type should "
5655 "never have an attribute created");
5657 if (RT == RegisterType::C) {
5658 if (Bindings.hasBindingInfoForDecl(VD))
5660 diag::warn_hlsl_user_defined_type_missing_member)
5661 <<
static_cast<int>(RT);
5669 if (DeclBindingInfo *BI = Bindings.getDeclBindingInfo(VD, RC)) {
5674 diag::warn_hlsl_user_defined_type_missing_member)
5675 <<
static_cast<int>(RT);
5683class InitListTransformer {
5687 QualType *DstIt =
nullptr;
5688 Expr **ArgIt =
nullptr;
5694 bool castInitializer(Expr *E) {
5695 assert(DstIt &&
"This should always be something!");
5696 if (DstIt == DestTypes.end()) {
5698 ArgExprs.push_back(E);
5703 DstIt = DestTypes.begin();
5706 Ctx, *DstIt,
false);
5711 ArgExprs.push_back(
Init);
5716 bool buildInitializerListImpl(Expr *E) {
5718 if (
auto *
Init = dyn_cast<InitListExpr>(E)) {
5719 for (
auto *SubInit :
Init->inits())
5720 if (!buildInitializerListImpl(SubInit))
5730 return castInitializer(E);
5732 if (
auto *VecTy = Ty->
getAs<VectorType>()) {
5737 for (uint64_t I = 0; I <
Size; ++I) {
5739 SizeTy, SourceLocation());
5745 if (!castInitializer(ElExpr.
get()))
5750 if (
auto *MTy = Ty->
getAs<ConstantMatrixType>()) {
5751 unsigned Rows = MTy->getNumRows();
5752 unsigned Cols = MTy->getNumColumns();
5753 QualType ElemTy = MTy->getElementType();
5755 for (
unsigned R = 0;
R < Rows; ++
R) {
5756 for (
unsigned C = 0;
C < Cols; ++
C) {
5769 if (!castInitializer(ElExpr.
get()))
5777 if (
auto *ArrTy = dyn_cast<ConstantArrayType>(Ty.
getTypePtr())) {
5781 for (uint64_t I = 0; I <
Size; ++I) {
5783 SizeTy, SourceLocation());
5788 if (!buildInitializerListImpl(ElExpr.
get()))
5795 llvm::SmallVector<CXXRecordDecl *> RecordDecls;
5796 RecordDecls.push_back(RD);
5801 MaterializeTemporaryExpr(Ty, E,
false);
5802 while (RecordDecls.back()->getNumBases()) {
5803 CXXRecordDecl *D = RecordDecls.back();
5805 "HLSL doesn't support multiple inheritance");
5806 RecordDecls.push_back(
5809 while (!RecordDecls.empty()) {
5810 CXXRecordDecl *RD = RecordDecls.pop_back_val();
5811 for (
auto *FD : RD->
fields()) {
5812 if (FD->isUnnamedBitField())
5820 if (!buildInitializerListImpl(Res.
get()))
5828 Expr *generateInitListsImpl(QualType Ty) {
5830 assert(ArgIt != ArgExprs.end() &&
"Something is off in iteration!");
5835 llvm::SmallVector<Expr *>
Inits;
5840 if (
auto *ATy = Ty->
getAs<VectorType>()) {
5841 ElTy = ATy->getElementType();
5842 Size = ATy->getNumElements();
5843 }
else if (
auto *CMTy = Ty->
getAs<ConstantMatrixType>()) {
5844 ElTy = CMTy->getElementType();
5845 Size = CMTy->getNumElementsFlattened();
5848 ElTy = VTy->getElementType();
5849 Size = VTy->getZExtSize();
5851 for (uint64_t I = 0; I <
Size; ++I)
5852 Inits.push_back(generateInitListsImpl(ElTy));
5855 llvm::SmallVector<CXXRecordDecl *> RecordDecls;
5856 RecordDecls.push_back(RD);
5857 while (RecordDecls.back()->getNumBases()) {
5858 CXXRecordDecl *D = RecordDecls.back();
5860 "HLSL doesn't support multiple inheritance");
5861 RecordDecls.push_back(
5864 while (!RecordDecls.empty()) {
5865 CXXRecordDecl *RD = RecordDecls.pop_back_val();
5866 for (
auto *FD : RD->
fields())
5867 if (!FD->isUnnamedBitField())
5871 auto *NewInit =
new (Ctx) InitListExpr(Ctx,
Inits.front()->getBeginLoc(),
5873 NewInit->setType(Ty);
5878 llvm::SmallVector<QualType, 16> DestTypes;
5879 llvm::SmallVector<Expr *, 16> ArgExprs;
5880 InitListTransformer(Sema &SemaRef,
const InitializedEntity &Entity)
5881 : S(SemaRef), Ctx(SemaRef.getASTContext()),
5882 Wrap(Entity.
getType()->isIncompleteArrayType()) {
5883 InitTy = Entity.
getType().getNonReferenceType();
5893 DstIt = DestTypes.begin();
5896 bool buildInitializerList(Expr *E) {
return buildInitializerListImpl(E); }
5898 Expr *generateInitLists() {
5899 assert(!ArgExprs.empty() &&
5900 "Call buildInitializerList to generate argument expressions.");
5901 ArgIt = ArgExprs.begin();
5903 return generateInitListsImpl(InitTy);
5904 llvm::SmallVector<Expr *>
Inits;
5905 while (ArgIt != ArgExprs.end())
5906 Inits.push_back(generateInitListsImpl(InitTy));
5908 auto *NewInit =
new (Ctx) InitListExpr(Ctx,
Inits.front()->getBeginLoc(),
5910 llvm::APInt ArySize(64,
Inits.size());
5912 ArraySizeModifier::Normal, 0));
5924 if (
const ArrayType *AT = dyn_cast<ArrayType>(Ty)) {
5931 if (
const auto *RT = Ty->
getAs<RecordType>()) {
5935 if (
const auto *CXXRD = dyn_cast<CXXRecordDecl>(RD)) {
5955 if (
Init->getType()->isScalarType())
5958 InitListTransformer ILT(
SemaRef, Entity);
5960 for (
unsigned I = 0; I <
Init->getNumInits(); ++I) {
5968 Init->setInit(I, E);
5970 if (!ILT.buildInitializerList(E))
5973 size_t ExpectedSize = ILT.DestTypes.size();
5974 size_t ActualSize = ILT.ArgExprs.size();
5975 if (ExpectedSize == 0 && ActualSize == 0)
5982 InitTy =
SemaRef.getASTContext().removeAddrSpaceQualType(InitTy);
5984 SemaRef.Diag(
Init->getBeginLoc(), diag::err_hlsl_incorrect_num_initializers)
5985 << (int)(ExpectedSize < ActualSize) << InitTy
5986 << ExpectedSize << ActualSize;
5996 assert(ExpectedSize > 0 &&
5997 "The expected size of an incomplete array type must be at least 1.");
5999 ((ActualSize + ExpectedSize - 1) / ExpectedSize) * ExpectedSize;
6007 InitTy =
SemaRef.getASTContext().removeAddrSpaceQualType(InitTy);
6008 if (ExpectedSize != ActualSize) {
6009 int TooManyOrFew = ActualSize > ExpectedSize ? 1 : 0;
6010 SemaRef.Diag(
Init->getBeginLoc(), diag::err_hlsl_incorrect_num_initializers)
6011 << TooManyOrFew << InitTy << ExpectedSize << ActualSize;
6018 Init->resizeInits(Ctx, NewInit->getNumInits());
6019 for (
unsigned I = 0; I < NewInit->getNumInits(); ++I)
6020 Init->updateInit(Ctx, I, NewInit->getInit(I));
6028 S.
Diag(OpLoc, diag::err_builtin_matrix_invalid_member)
6038 StringRef AccessorName = CompName->
getName();
6039 assert(!AccessorName.empty() &&
"Matrix Accessor must have a name");
6041 unsigned Rows = MT->getNumRows();
6042 unsigned Cols = MT->getNumColumns();
6043 bool IsZeroBasedAccessor =
false;
6044 unsigned ChunkLen = 0;
6045 if (AccessorName.size() < 2)
6047 "length 4 for zero based: \'_mRC\' or "
6048 "length 3 for one-based: \'_RC\' accessor",
6051 if (AccessorName[0] ==
'_') {
6052 if (AccessorName[1] ==
'm') {
6053 IsZeroBasedAccessor =
true;
6060 S, AccessorName,
"zero based: \'_mRC\' or one-based: \'_RC\' accessor",
6063 if (AccessorName.size() % ChunkLen != 0) {
6064 const llvm::StringRef
Expected = IsZeroBasedAccessor
6065 ?
"zero based: '_mRC' accessor"
6066 :
"one-based: '_RC' accessor";
6071 auto isDigit = [](
char c) {
return c >=
'0' &&
c <=
'9'; };
6072 auto isZeroBasedIndex = [](
unsigned i) {
return i <= 3; };
6073 auto isOneBasedIndex = [](
unsigned i) {
return i >= 1 && i <= 4; };
6075 bool HasRepeated =
false;
6077 unsigned NumComponents = 0;
6078 const char *Begin = AccessorName.data();
6080 for (
unsigned I = 0, E = AccessorName.size(); I < E; I += ChunkLen) {
6081 const char *Chunk = Begin + I;
6082 char RowChar = 0, ColChar = 0;
6083 if (IsZeroBasedAccessor) {
6085 if (Chunk[0] !=
'_' || Chunk[1] !=
'm') {
6086 char Bad = (Chunk[0] !=
'_') ? Chunk[0] : Chunk[1];
6088 S, StringRef(&Bad, 1),
"\'_m\' prefix",
6095 if (Chunk[0] !=
'_')
6097 S, StringRef(&Chunk[0], 1),
"\'_\' prefix",
6104 bool IsDigitsError =
false;
6106 unsigned BadPos = IsZeroBasedAccessor ? 2 : 1;
6110 IsDigitsError =
true;
6114 unsigned BadPos = IsZeroBasedAccessor ? 3 : 2;
6118 IsDigitsError =
true;
6123 unsigned Row = RowChar -
'0';
6124 unsigned Col = ColChar -
'0';
6126 bool HasIndexingError =
false;
6127 if (IsZeroBasedAccessor) {
6129 if (!isZeroBasedIndex(Row)) {
6130 S.
Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
6132 HasIndexingError =
true;
6134 if (!isZeroBasedIndex(Col)) {
6135 S.
Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
6137 HasIndexingError =
true;
6141 if (!isOneBasedIndex(Row)) {
6142 S.
Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
6144 HasIndexingError =
true;
6146 if (!isOneBasedIndex(Col)) {
6147 S.
Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
6149 HasIndexingError =
true;
6156 if (HasIndexingError)
6162 bool HasBoundsError =
false;
6164 Diag(OpLoc, diag::err_hlsl_matrix_index_out_of_bounds)
6166 HasBoundsError =
true;
6169 Diag(OpLoc, diag::err_hlsl_matrix_index_out_of_bounds)
6171 HasBoundsError =
true;
6176 unsigned FlatIndex = Row * Cols + Col;
6177 if (Seen[FlatIndex])
6179 Seen[FlatIndex] =
true;
6182 if (NumComponents == 0 || NumComponents > 4) {
6183 S.
Diag(OpLoc, diag::err_hlsl_matrix_swizzle_invalid_length)
6188 QualType ElemTy = MT->getElementType();
6189 if (NumComponents == 1)
6195 for (Sema::ExtVectorDeclsType::iterator
6199 if ((*I)->getUnderlyingType() == VT)
6210 trackLocalResource(VDecl,
Init);
6212 const HLSLVkConstantIdAttr *ConstIdAttr =
6213 VDecl->
getAttr<HLSLVkConstantIdAttr>();
6220 if (!
Init->isCXX11ConstantExpr(Context, &InitValue)) {
6230 int ConstantID = ConstIdAttr->getId();
6231 llvm::APInt IDVal(Context.getIntWidth(Context.IntTy), ConstantID);
6233 ConstIdAttr->getLocation());
6237 if (
C->getType()->getCanonicalTypeUnqualified() !=
6241 Context.getTrivialTypeSourceInfo(
6242 Init->getType(),
Init->getExprLoc()),
6261 if (!Params || Params->
size() != 1)
6274 if (
auto *TTP = dyn_cast<TemplateTypeParmDecl>(P)) {
6275 if (TTP->hasDefaultArgument()) {
6276 TemplateArgs.
addArgument(TTP->getDefaultArgument());
6279 }
else if (
auto *NTTP = dyn_cast<NonTypeTemplateParmDecl>(P)) {
6280 if (NTTP->hasDefaultArgument()) {
6281 TemplateArgs.
addArgument(NTTP->getDefaultArgument());
6284 }
else if (
auto *TTPD = dyn_cast<TemplateTemplateParmDecl>(P)) {
6285 if (TTPD->hasDefaultArgument()) {
6286 TemplateArgs.
addArgument(TTPD->getDefaultArgument());
6293 return SemaRef.CheckTemplateIdType(
6295 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 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 bool hasCounterHandle(const CXXRecordDecl *RD)
static bool CheckVectorSelect(Sema *S, CallExpr *TheCall)
static QualType handleFloatVectorBinOpConversion(Sema &SemaRef, ExprResult &LHS, ExprResult &RHS, QualType LHSType, QualType RHSType, QualType LElTy, QualType RElTy, bool IsCompAssign)
static ResourceClass getResourceClass(RegisterType RT)
static CXXRecordDecl * createHostLayoutStruct(Sema &S, CXXRecordDecl *StructDecl)
static bool CheckScalarOrVector(Sema *S, CallExpr *TheCall, QualType Scalar, unsigned ArgIndex)
static bool CheckSamplingBuiltin(Sema &S, CallExpr *TheCall, SampleKind Kind)
static bool CheckScalarOrVectorOrMatrix(Sema *S, CallExpr *TheCall, QualType Scalar, unsigned ArgIndex)
static bool CheckFloatRepresentation(Sema *S, SourceLocation Loc, int ArgOrdinal, clang::QualType PassedType)
static bool CheckAnyDoubleRepresentation(Sema *S, SourceLocation Loc, int ArgOrdinal, clang::QualType PassedType)
static bool requiresImplicitBufferLayoutStructure(const CXXRecordDecl *RD)
static bool CheckResourceHandle(Sema *S, CallExpr *TheCall, unsigned ArgIndex, llvm::function_ref< bool(const HLSLAttributedResourceType *ResType)> Check=nullptr)
static void validatePackoffset(Sema &S, HLSLBufferDecl *BufDecl)
static bool IsDefaultBufferConstantDecl(const ASTContext &Ctx, VarDecl *VD)
HLSLResourceBindingAttr::RegisterType RegisterType
static CastKind getScalarCastKind(ASTContext &Ctx, QualType DestTy, QualType SrcTy)
static bool CheckGatherBuiltin(Sema &S, CallExpr *TheCall, bool IsCmp)
static bool isValidWaveSizeValue(unsigned Value)
static bool isResourceRecordTypeOrArrayOf(QualType Ty)
static bool AccumulateHLSLResourceSlots(QualType Ty, uint64_t &StartSlot, const uint64_t &Limit, const ResourceClass ResClass, ASTContext &Ctx, uint64_t ArrayCount=1)
static bool CheckNoDoubleVectors(Sema *S, SourceLocation Loc, int ArgOrdinal, clang::QualType PassedType)
static bool ValidateMultipleRegisterAnnotations(Sema &S, Decl *TheDecl, RegisterType regType)
static bool CheckTextureSamplerAndLocation(Sema &S, CallExpr *TheCall)
static bool DiagnoseLocalRegisterBinding(Sema &S, SourceLocation &ArgLoc, Decl *D, RegisterType RegType, bool SpecifiedSpace)
static bool CheckIndexType(Sema *S, CallExpr *TheCall, unsigned IndexArgIndex)
This file declares semantic analysis for HLSL constructs.
Defines the clang::SourceLocation class and associated facilities.
Defines various enumerations that describe declaration and type specifiers.
C Language Family Type Representation.
Defines the clang::TypeLoc interface and its subclasses.
C Language Family Type Representation.
static const TypeInfo & getInfo(unsigned id)
__device__ __2f16 float c
return(__x > > __y)|(__x<<(32 - __y))
APValue - This class implements a discriminated union of [uninitialized] [APSInt] [APFloat],...
virtual bool HandleTopLevelDecl(DeclGroupRef D)
HandleTopLevelDecl - Handle the specified top-level declaration.
Holds long-lived AST nodes (such as types and decls) that can be referred to throughout the semantic ...
unsigned getIntWidth(QualType T) const
int getIntegerTypeOrder(QualType LHS, QualType RHS) const
Return the highest ranked integer type, see C99 6.3.1.8p1.
QualType getPointerType(QualType T) const
Return the uniqued reference to the type for a pointer to the specified type.
const IncompleteArrayType * getAsIncompleteArrayType(QualType T) const
QualType getConstantArrayType(QualType EltTy, const llvm::APInt &ArySize, const Expr *SizeExpr, ArraySizeModifier ASM, unsigned IndexTypeQuals) const
Return the unique reference to the type for a constant array of the specified element type.
QualType getBaseElementType(const ArrayType *VAT) const
Return the innermost element type of an array type.
int getFloatingTypeOrder(QualType LHS, QualType RHS) const
Compare the rank of the two specified floating point types, ignoring the domain of the type (i....
TypeSourceInfo * getTrivialTypeSourceInfo(QualType T, SourceLocation Loc=SourceLocation()) const
Allocate a TypeSourceInfo where all locations have been initialized to a given location,...
QualType getStringLiteralArrayType(QualType EltTy, unsigned Length) const
Return a type for a constant array for a string literal of the specified element type and length.
uint64_t getTypeSize(QualType T) const
Return the size of the specified (complete) type T, in bits.
CharUnits getTypeSizeInChars(QualType T) const
Return the size of the specified (complete) type T, in characters.
CanQualType UnsignedIntTy
QualType getTypedefType(ElaboratedTypeKeyword Keyword, NestedNameSpecifier Qualifier, const TypedefNameDecl *Decl, QualType UnderlyingType=QualType(), std::optional< bool > TypeMatchesDeclOrNone=std::nullopt) const
Return the unique reference to the type for the specified typedef-name decl.
llvm::StringRef backupStr(llvm::StringRef S) const
QualType getSizeType() const
Return the unique type for "size_t" (C99 7.17), defined in <stddef.h>.
QualType getExtVectorType(QualType VectorType, unsigned NumElts) const
Return the unique reference to an extended vector type of the specified element type and size.
const TargetInfo & getTargetInfo() const
QualType getHLSLAttributedResourceType(QualType Wrapped, QualType Contained, const HLSLAttributedResourceType::Attributes &Attrs)
QualType getAddrSpaceQualType(QualType T, LangAS AddressSpace) const
Return the uniqued reference to the type for an address space qualified type with the specified type ...
CanQualType getCanonicalTagType(const TagDecl *TD) const
static bool hasSameUnqualifiedType(QualType T1, QualType T2)
Determine whether the given types are equivalent after cvr-qualifiers have been removed.
QualType getConstantMatrixType(QualType ElementType, unsigned NumRows, unsigned NumColumns) const
Return the unique reference to the matrix type of the specified element type and size.
unsigned getTypeAlign(QualType T) const
Return the ABI-specified alignment of a (complete) type T, in bits.
Represents an array type, per C99 6.7.5.2 - Array Declarators.
QualType getElementType() const
Attr - This represents one attribute.
attr::Kind getKind() const
SourceLocation getLocation() const
SourceLocation getScopeLoc() const
SourceRange getRange() const
const IdentifierInfo * getScopeName() const
SourceLocation getLoc() const
const IdentifierInfo * getAttrName() const
Represents a base class of a C++ class.
QualType getType() const
Retrieves the type of the base class.
Represents a static or instance method of a struct/union/class.
Represents a C++ struct/union/class.
bool isHLSLIntangible() const
Returns true if the class contains HLSL intangible type, either as a field or in base class.
static CXXRecordDecl * Create(const ASTContext &C, TagKind TK, DeclContext *DC, SourceLocation StartLoc, SourceLocation IdLoc, IdentifierInfo *Id, CXXRecordDecl *PrevDecl=nullptr)
void setBases(CXXBaseSpecifier const *const *Bases, unsigned NumBases)
Sets the base classes of this struct or class.
base_class_iterator bases_end()
void completeDefinition() override
Indicates that the definition of this class is now complete.
unsigned getNumBases() const
Retrieves the number of base classes of this class.
base_class_iterator bases_begin()
bool isEmpty() const
Determine whether this is an empty class in the sense of (C++11 [meta.unary.prop]).
CallExpr - Represents a function call (C99 6.5.2.2, C++ [expr.call]).
Expr * getArg(unsigned Arg)
getArg - Return the specified argument.
SourceLocation getBeginLoc() const
static CallExpr * Create(const ASTContext &Ctx, Expr *Fn, ArrayRef< Expr * > Args, QualType Ty, ExprValueKind VK, SourceLocation RParenLoc, FPOptionsOverride FPFeatures, unsigned MinNumArgs=0, ADLCallKind UsesADL=NotADL)
Create a call expression.
FunctionDecl * getDirectCallee()
If the callee is a FunctionDecl, return it. Otherwise return null.
unsigned getNumArgs() const
getNumArgs - Return the number of actual arguments to this call.
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...
void setValueKind(ExprValueKind Cat)
setValueKind - Set the value kind produced by this expression.
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.
IdentifierInfo * getIdentifier() const
Get the identifier that names this declaration, if there is one.
StringRef getName() const
Get the name of identifier for this declaration as a StringRef.
DeclarationName getDeclName() const
Get the actual, stored name of the declaration, which may be a special name.
A C++ nested-name-specifier augmented with source location information.
OpaqueValueExpr - An expression referring to an opaque object of a fixed type and value class.
Represents a parameter to a function.
ParsedAttr - Represents a syntactic attribute.
unsigned getSemanticSpelling() const
If the parsed attribute has a semantic equivalent, and it would have a semantic Spelling enumeration ...
unsigned getMinArgs() const
bool checkExactlyNumArgs(class Sema &S, unsigned Num) const
Check if the attribute has exactly as many args as Num.
IdentifierLoc * getArgAsIdent(unsigned Arg) const
bool hasParsedType() const
const ParsedType & getTypeArg() const
unsigned getNumArgs() const
getNumArgs - Return the number of actual arguments to this attribute.
bool isArgIdent(unsigned Arg) const
Expr * getArgAsExpr(unsigned Arg) const
AttributeCommonInfo::Kind getKind() const
A (possibly-)qualified type.
void addRestrict()
Add the restrict qualifier to this QualType.
QualType getNonLValueExprType(const ASTContext &Context) const
Determine the type of a (typically non-lvalue) expression with the specified result type.
QualType getDesugaredType(const ASTContext &Context) const
Return the specified type with any "sugar" removed from the type.
bool isNull() const
Return true if this QualType doesn't point to a type yet.
const Type * getTypePtr() const
Retrieves a pointer to the underlying (unqualified) type.
LangAS getAddressSpace() const
Return the address space of this type.
QualType getNonReferenceType() const
If Type is a reference type (e.g., const int&), returns the type that the reference refers to ("const...
QualType getCanonicalType() const
QualType getUnqualifiedType() const
Retrieve the unqualified variant of the given type, removing as little sugar as possible.
bool hasAddressSpace() const
Check if this type has any address space qualifier.
Represents a struct/union/class.
field_iterator field_end() const
field_range fields() const
field_iterator field_begin() 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)
bool diagnosePositionType(QualType T, const ParsedAttr &AL)
bool handleInitialization(VarDecl *VDecl, Expr *&Init)
bool diagnoseInputIDType(QualType T, const ParsedAttr &AL)
void handleParamModifierAttr(Decl *D, const ParsedAttr &AL)
bool CheckResourceBinOp(BinaryOperatorKind Opc, Expr *LHSExpr, Expr *RHSExpr, SourceLocation Loc)
bool CanPerformAggregateSplatCast(Expr *Src, QualType DestType)
bool ActOnResourceMemberAccessExpr(MemberExpr *ME)
bool IsScalarizedLayoutCompatible(QualType T1, QualType T2) const
QualType ActOnTemplateShorthand(TemplateDecl *Template, SourceLocation NameLoc)
void diagnoseSystemSemanticAttr(Decl *D, const ParsedAttr &AL, std::optional< unsigned > Index)
void handleRootSignatureAttr(Decl *D, const ParsedAttr &AL)
bool CheckCompatibleParameterABI(FunctionDecl *New, FunctionDecl *Old)
QualType handleVectorBinOpConversion(ExprResult &LHS, ExprResult &RHS, QualType LHSType, QualType RHSType, bool IsCompAssign)
QualType checkMatrixComponent(Sema &S, QualType baseType, ExprValueKind &VK, SourceLocation OpLoc, const IdentifierInfo *CompName, SourceLocation CompLoc)
void handleResourceBindingAttr(Decl *D, const ParsedAttr &AL)
bool IsTypedResourceElementCompatible(QualType T1)
bool transformInitList(const InitializedEntity &Entity, InitListExpr *Init)
void handleNumThreadsAttr(Decl *D, const ParsedAttr &AL)
bool ActOnUninitializedVarDecl(VarDecl *D)
void handleVkExtBuiltinInputAttr(Decl *D, const ParsedAttr &AL)
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)
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)
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 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 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)
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