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);
1495 bool HadError =
false;
1496 auto ReportError = [
this, &HadError](
SourceLocation Loc, uint32_t LowerBound,
1497 uint32_t UpperBound) {
1499 this->
Diag(Loc, diag::err_hlsl_invalid_rootsig_value)
1500 << LowerBound << UpperBound;
1507 this->
Diag(Loc, diag::err_hlsl_invalid_rootsig_value)
1508 << llvm::formatv(
"{0:f}", LowerBound).sstr<6>()
1509 << llvm::formatv(
"{0:f}", UpperBound).sstr<6>();
1512 auto VerifyRegister = [ReportError](
SourceLocation Loc, uint32_t Register) {
1513 if (!llvm::hlsl::rootsig::verifyRegisterValue(Register))
1514 ReportError(Loc, 0, 0xfffffffe);
1517 auto VerifySpace = [ReportError](
SourceLocation Loc, uint32_t Space) {
1518 if (!llvm::hlsl::rootsig::verifyRegisterSpace(Space))
1519 ReportError(Loc, 0, 0xffffffef);
1522 const uint32_t Version =
1523 llvm::to_underlying(
SemaRef.getLangOpts().HLSLRootSigVer);
1524 const uint32_t VersionEnum = Version - 1;
1525 auto ReportFlagError = [
this, &HadError, VersionEnum](
SourceLocation Loc) {
1527 this->
Diag(Loc, diag::err_hlsl_invalid_rootsig_flag)
1534 const llvm::hlsl::rootsig::RootElement &Elem = RootSigElem.
getElement();
1535 if (
const auto *Descriptor =
1536 std::get_if<llvm::hlsl::rootsig::RootDescriptor>(&Elem)) {
1537 VerifyRegister(Loc, Descriptor->Reg.Number);
1538 VerifySpace(Loc, Descriptor->Space);
1540 if (!llvm::hlsl::rootsig::verifyRootDescriptorFlag(Version,
1542 ReportFlagError(Loc);
1543 }
else if (
const auto *Constants =
1544 std::get_if<llvm::hlsl::rootsig::RootConstants>(&Elem)) {
1545 VerifyRegister(Loc, Constants->Reg.Number);
1546 VerifySpace(Loc, Constants->Space);
1547 }
else if (
const auto *Sampler =
1548 std::get_if<llvm::hlsl::rootsig::StaticSampler>(&Elem)) {
1549 VerifyRegister(Loc, Sampler->Reg.Number);
1550 VerifySpace(Loc, Sampler->Space);
1553 "By construction, parseFloatParam can't produce a NaN from a "
1554 "float_literal token");
1556 if (!llvm::hlsl::rootsig::verifyMaxAnisotropy(Sampler->MaxAnisotropy))
1557 ReportError(Loc, 0, 16);
1558 if (!llvm::hlsl::rootsig::verifyMipLODBias(Sampler->MipLODBias))
1559 ReportFloatError(Loc, -16.f, 15.99f);
1560 }
else if (
const auto *Clause =
1561 std::get_if<llvm::hlsl::rootsig::DescriptorTableClause>(
1563 VerifyRegister(Loc, Clause->Reg.Number);
1564 VerifySpace(Loc, Clause->Space);
1566 if (!llvm::hlsl::rootsig::verifyNumDescriptors(Clause->NumDescriptors)) {
1570 ReportError(Loc, 1, 0xfffffffe);
1573 if (!llvm::hlsl::rootsig::verifyDescriptorRangeFlag(Version, Clause->Type,
1575 ReportFlagError(Loc);
1579 PerVisibilityBindingChecker BindingChecker(
this);
1580 SmallVector<std::pair<
const llvm::hlsl::rootsig::DescriptorTableClause *,
1585 const llvm::hlsl::rootsig::RootElement &Elem = RootSigElem.
getElement();
1586 if (
const auto *Descriptor =
1587 std::get_if<llvm::hlsl::rootsig::RootDescriptor>(&Elem)) {
1588 uint32_t LowerBound(Descriptor->Reg.Number);
1589 uint32_t UpperBound(LowerBound);
1591 BindingChecker.trackBinding(
1592 Descriptor->Visibility,
1593 static_cast<llvm::dxil::ResourceClass
>(Descriptor->Type),
1594 Descriptor->Space, LowerBound, UpperBound, &RootSigElem);
1595 }
else if (
const auto *Constants =
1596 std::get_if<llvm::hlsl::rootsig::RootConstants>(&Elem)) {
1597 uint32_t LowerBound(Constants->Reg.Number);
1598 uint32_t UpperBound(LowerBound);
1600 BindingChecker.trackBinding(
1601 Constants->Visibility, llvm::dxil::ResourceClass::CBuffer,
1602 Constants->Space, LowerBound, UpperBound, &RootSigElem);
1603 }
else if (
const auto *Sampler =
1604 std::get_if<llvm::hlsl::rootsig::StaticSampler>(&Elem)) {
1605 uint32_t LowerBound(Sampler->Reg.Number);
1606 uint32_t UpperBound(LowerBound);
1608 BindingChecker.trackBinding(
1609 Sampler->Visibility, llvm::dxil::ResourceClass::Sampler,
1610 Sampler->Space, LowerBound, UpperBound, &RootSigElem);
1611 }
else if (
const auto *Clause =
1612 std::get_if<llvm::hlsl::rootsig::DescriptorTableClause>(
1615 UnboundClauses.emplace_back(Clause, &RootSigElem);
1616 }
else if (
const auto *Table =
1617 std::get_if<llvm::hlsl::rootsig::DescriptorTable>(&Elem)) {
1618 assert(UnboundClauses.size() == Table->NumClauses &&
1619 "Number of unbound elements must match the number of clauses");
1620 bool HasAnySampler =
false;
1621 bool HasAnyNonSampler =
false;
1622 uint64_t Offset = 0;
1623 bool IsPrevUnbound =
false;
1624 for (
const auto &[Clause, ClauseElem] : UnboundClauses) {
1626 if (Clause->Type == llvm::dxil::ResourceClass::Sampler)
1627 HasAnySampler =
true;
1629 HasAnyNonSampler =
true;
1631 if (HasAnySampler && HasAnyNonSampler)
1632 Diag(Loc, diag::err_hlsl_invalid_mixed_resources);
1637 if (Clause->NumDescriptors == 0)
1641 Clause->Offset == llvm::hlsl::rootsig::DescriptorTableOffsetAppend;
1643 Offset = Clause->Offset;
1645 uint64_t RangeBound = llvm::hlsl::rootsig::computeRangeBound(
1646 Offset, Clause->NumDescriptors);
1648 if (IsPrevUnbound && IsAppending)
1649 Diag(Loc, diag::err_hlsl_appending_onto_unbound);
1650 else if (!llvm::hlsl::rootsig::verifyNoOverflowedOffset(RangeBound))
1651 Diag(Loc, diag::err_hlsl_offset_overflow) << Offset << RangeBound;
1654 Offset = RangeBound + 1;
1655 IsPrevUnbound = Clause->NumDescriptors ==
1656 llvm::hlsl::rootsig::NumDescriptorsUnbounded;
1659 uint32_t LowerBound(Clause->Reg.Number);
1660 uint32_t UpperBound = llvm::hlsl::rootsig::computeRangeBound(
1661 LowerBound, Clause->NumDescriptors);
1663 BindingChecker.trackBinding(
1665 static_cast<llvm::dxil::ResourceClass
>(Clause->Type), Clause->Space,
1666 LowerBound, UpperBound, ClauseElem);
1668 UnboundClauses.clear();
1672 return BindingChecker.checkOverlap();
1677 Diag(AL.
getLoc(), diag::err_attribute_wrong_number_arguments) << AL << 1;
1682 if (
auto *RS = D->
getAttr<RootSignatureAttr>()) {
1683 if (RS->getSignatureIdent() != Ident) {
1684 Diag(AL.
getLoc(), diag::err_disallowed_duplicate_attribute) << RS;
1688 Diag(AL.
getLoc(), diag::warn_duplicate_attribute_exact) << RS;
1694 if (
auto *SignatureDecl =
1695 dyn_cast<HLSLRootSignatureDecl>(R.getFoundDecl())) {
1702 llvm::VersionTuple SMVersion =
1707 uint32_t ZMax = 1024;
1708 uint32_t ThreadMax = 1024;
1709 if (IsDXIL && SMVersion.getMajor() <= 4) {
1712 }
else if (IsDXIL && SMVersion.getMajor() == 5) {
1722 diag::err_hlsl_numthreads_argument_oor)
1731 diag::err_hlsl_numthreads_argument_oor)
1740 diag::err_hlsl_numthreads_argument_oor)
1745 if (
X * Y * Z > ThreadMax) {
1746 Diag(AL.
getLoc(), diag::err_hlsl_numthreads_invalid) << ThreadMax;
1763 if (SpelledArgsCount == 0 || SpelledArgsCount > 3)
1771 if (SpelledArgsCount > 1 &&
1775 uint32_t Preferred = 0;
1776 if (SpelledArgsCount > 2 &&
1780 if (SpelledArgsCount > 2) {
1783 diag::err_attribute_power_of_two_in_range)
1784 << AL << llvm::dxil::MinWaveSize << llvm::dxil::MaxWaveSize
1789 if (Preferred < Min || Preferred >
Max) {
1791 diag::err_attribute_power_of_two_in_range)
1792 << AL <<
Min <<
Max << Preferred;
1795 }
else if (SpelledArgsCount > 1) {
1798 diag::err_attribute_power_of_two_in_range)
1799 << AL << llvm::dxil::MinWaveSize << llvm::dxil::MaxWaveSize <<
Max;
1803 Diag(AL.
getLoc(), diag::err_attribute_argument_invalid) << AL << 1;
1806 Diag(AL.
getLoc(), diag::warn_attr_min_eq_max) << AL;
1811 diag::err_attribute_power_of_two_in_range)
1812 << AL << llvm::dxil::MinWaveSize << llvm::dxil::MaxWaveSize <<
Min;
1817 HLSLWaveSizeAttr *NewAttr =
1854 uint32_t Binding = 0;
1878 if (!T->hasUnsignedIntegerRepresentation() ||
1879 (VT && VT->getNumElements() > 3)) {
1880 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type)
1881 << AL <<
"uint/uint2/uint3";
1890 if (!T->hasFloatingRepresentation() || (VT && VT->getNumElements() > 4)) {
1891 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type)
1892 << AL <<
"float/float1/float2/float3/float4";
1900 std::optional<unsigned> Index) {
1904 QualType ValueType = VD->getType();
1905 if (
auto *FD = dyn_cast<FunctionDecl>(D))
1908 bool IsOutput =
false;
1909 if (HLSLParamModifierAttr *MA = D->
getAttr<HLSLParamModifierAttr>()) {
1916 if (SemanticName ==
"SV_DISPATCHTHREADID") {
1919 Diag(AL.
getLoc(), diag::err_hlsl_semantic_output_not_supported) << AL;
1920 if (Index.has_value())
1921 Diag(AL.
getLoc(), diag::err_hlsl_semantic_indexing_not_supported) << AL;
1926 if (SemanticName ==
"SV_GROUPINDEX") {
1928 Diag(AL.
getLoc(), diag::err_hlsl_semantic_output_not_supported) << AL;
1929 if (Index.has_value())
1930 Diag(AL.
getLoc(), diag::err_hlsl_semantic_indexing_not_supported) << AL;
1935 if (SemanticName ==
"SV_GROUPTHREADID") {
1938 Diag(AL.
getLoc(), diag::err_hlsl_semantic_output_not_supported) << AL;
1939 if (Index.has_value())
1940 Diag(AL.
getLoc(), diag::err_hlsl_semantic_indexing_not_supported) << AL;
1945 if (SemanticName ==
"SV_GROUPID") {
1948 Diag(AL.
getLoc(), diag::err_hlsl_semantic_output_not_supported) << AL;
1949 if (Index.has_value())
1950 Diag(AL.
getLoc(), diag::err_hlsl_semantic_indexing_not_supported) << AL;
1955 if (SemanticName ==
"SV_POSITION") {
1956 const auto *VT = ValueType->getAs<
VectorType>();
1957 if (!ValueType->hasFloatingRepresentation() ||
1958 (VT && VT->getNumElements() > 4))
1959 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type)
1960 << AL <<
"float/float1/float2/float3/float4";
1965 if (SemanticName ==
"SV_VERTEXID") {
1966 uint64_t SizeInBits =
SemaRef.Context.getTypeSize(ValueType);
1967 if (!ValueType->isUnsignedIntegerType() || SizeInBits != 32)
1968 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type) << AL <<
"uint";
1973 if (SemanticName ==
"SV_TARGET") {
1974 const auto *VT = ValueType->getAs<
VectorType>();
1975 if (!ValueType->hasFloatingRepresentation() ||
1976 (VT && VT->getNumElements() > 4))
1977 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type)
1978 << AL <<
"float/float1/float2/float3/float4";
1983 Diag(AL.
getLoc(), diag::err_hlsl_unknown_semantic) << AL;
1987 uint32_t IndexValue(0), ExplicitIndex(0);
1990 assert(0 &&
"HLSLUnparsedSemantic is expected to have 2 int arguments.");
1992 assert(IndexValue > 0 ? ExplicitIndex :
true);
1993 std::optional<unsigned> Index =
1994 ExplicitIndex ? std::optional<unsigned>(IndexValue) : std::nullopt;
2004 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_ast_node)
2005 << AL <<
"shader constant in a constant buffer";
2009 uint32_t SubComponent;
2019 bool IsAggregateTy = (T->isArrayType() || T->isStructureType());
2024 if (IsAggregateTy) {
2025 Diag(AL.
getLoc(), diag::err_hlsl_invalid_register_or_packoffset);
2029 if ((Component * 32 + Size) > 128) {
2030 Diag(AL.
getLoc(), diag::err_hlsl_packoffset_cross_reg_boundary);
2035 EltTy = VT->getElementType();
2037 if (Align > 32 && Component == 1) {
2040 Diag(AL.
getLoc(), diag::err_hlsl_packoffset_alignment_mismatch)
2054 if (!
SemaRef.checkStringLiteralArgumentAttr(AL, 0, Str, &ArgLoc))
2057 llvm::Triple::EnvironmentType ShaderType;
2058 if (!HLSLShaderAttr::ConvertStrToEnvironmentType(Str, ShaderType)) {
2059 Diag(AL.
getLoc(), diag::warn_attribute_type_not_supported)
2060 << AL << Str << ArgLoc;
2074 assert(AttrList.size() &&
"expected list of resource attributes");
2081 HLSLAttributedResourceType::Attributes ResAttrs;
2083 bool HasResourceClass =
false;
2084 bool HasResourceDimension =
false;
2085 for (
const Attr *A : AttrList) {
2090 case attr::HLSLResourceClass: {
2092 if (HasResourceClass) {
2094 ? diag::warn_duplicate_attribute_exact
2095 : diag::warn_duplicate_attribute)
2099 ResAttrs.ResourceClass = RC;
2100 HasResourceClass =
true;
2103 case attr::HLSLResourceDimension: {
2104 llvm::dxil::ResourceDimension RD =
2106 if (HasResourceDimension) {
2108 ? diag::warn_duplicate_attribute_exact
2109 : diag::warn_duplicate_attribute)
2113 ResAttrs.ResourceDimension = RD;
2114 HasResourceDimension =
true;
2118 if (ResAttrs.IsROV) {
2122 ResAttrs.IsROV =
true;
2124 case attr::HLSLRawBuffer:
2125 if (ResAttrs.RawBuffer) {
2129 ResAttrs.RawBuffer =
true;
2131 case attr::HLSLIsCounter:
2132 if (ResAttrs.IsCounter) {
2136 ResAttrs.IsCounter =
true;
2138 case attr::HLSLContainedType: {
2141 if (!ContainedTy.
isNull()) {
2143 ? diag::warn_duplicate_attribute_exact
2144 : diag::warn_duplicate_attribute)
2153 llvm_unreachable(
"unhandled resource attribute type");
2157 if (!HasResourceClass) {
2158 S.
Diag(AttrList.back()->getRange().getEnd(),
2159 diag::err_hlsl_missing_resource_class);
2164 Wrapped, ContainedTy, ResAttrs);
2166 if (LocInfo && ContainedTyInfo) {
2179 if (!T->isHLSLResourceType()) {
2180 Diag(AL.
getLoc(), diag::err_hlsl_attribute_needs_intangible_type)
2195 AttributeCommonInfo::AS_CXX11, 0, false ,
2200 case ParsedAttr::AT_HLSLResourceClass: {
2202 Diag(AL.
getLoc(), diag::err_attribute_argument_type)
2213 if (!HLSLResourceClassAttr::ConvertStrToResourceClass(Identifier, RC)) {
2214 Diag(ArgLoc, diag::warn_attribute_type_not_supported)
2215 <<
"ResourceClass" << Identifier;
2218 A = HLSLResourceClassAttr::Create(
getASTContext(), RC, ACI);
2222 case ParsedAttr::AT_HLSLResourceDimension: {
2223 StringRef Identifier;
2225 if (!
SemaRef.checkStringLiteralArgumentAttr(AL, 0, Identifier, &ArgLoc))
2229 llvm::dxil::ResourceDimension RD;
2230 if (!HLSLResourceDimensionAttr::ConvertStrToResourceDimension(Identifier,
2232 Diag(ArgLoc, diag::warn_attribute_type_not_supported)
2233 <<
"ResourceDimension" << Identifier;
2236 A = HLSLResourceDimensionAttr::Create(
getASTContext(), RD, ACI);
2240 case ParsedAttr::AT_HLSLROV:
2244 case ParsedAttr::AT_HLSLRawBuffer:
2248 case ParsedAttr::AT_HLSLIsCounter:
2252 case ParsedAttr::AT_HLSLContainedType: {
2254 Diag(AL.
getLoc(), diag::err_attribute_wrong_number_arguments) << AL << 1;
2260 assert(TSI &&
"no type source info for attribute argument");
2262 diag::err_incomplete_type))
2264 A = HLSLContainedTypeAttr::Create(
getASTContext(), TSI, ACI);
2269 llvm_unreachable(
"unhandled HLSL attribute");
2272 HLSLResourcesTypeAttrs.emplace_back(A);
2278 if (!HLSLResourcesTypeAttrs.size())
2284 HLSLResourcesTypeAttrs, QT, &LocInfo)) {
2285 const HLSLAttributedResourceType *RT =
2292 LocsForHLSLAttributedResources.insert(std::pair(RT, LocInfo));
2294 HLSLResourcesTypeAttrs.clear();
2302 auto I = LocsForHLSLAttributedResources.find(RT);
2303 if (I != LocsForHLSLAttributedResources.end()) {
2304 LocInfo = I->second;
2305 LocsForHLSLAttributedResources.erase(I);
2314void SemaHLSL::collectResourceBindingsOnUserRecordDecl(
const VarDecl *VD,
2315 const RecordType *RT) {
2316 const RecordDecl *RD = RT->getDecl()->getDefinitionOrSelf();
2323 "incomplete arrays inside user defined types are not supported");
2332 if (
const HLSLAttributedResourceType *AttrResType =
2333 HLSLAttributedResourceType::findHandleTypeOnResource(Ty)) {
2338 Bindings.addDeclBindingInfo(VD, RC);
2339 }
else if (
const RecordType *RT = dyn_cast<RecordType>(Ty)) {
2345 collectResourceBindingsOnUserRecordDecl(VD, RT);
2357 bool SpecifiedSpace) {
2358 int RegTypeNum =
static_cast<int>(RegType);
2361 if (D->
hasAttr<HLSLGroupSharedAddressSpaceAttr>()) {
2362 S.
Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
2367 if (
HLSLBufferDecl *CBufferOrTBuffer = dyn_cast<HLSLBufferDecl>(D)) {
2368 ResourceClass RC = CBufferOrTBuffer->isCBuffer() ? ResourceClass::CBuffer
2369 : ResourceClass::SRV;
2379 assert(
isa<VarDecl>(D) &&
"D is expected to be VarDecl or HLSLBufferDecl");
2383 if (
const HLSLAttributedResourceType *AttrResType =
2384 HLSLAttributedResourceType::findHandleTypeOnResource(
2401 if (SpecifiedSpace && !DeclaredInCOrTBuffer)
2402 S.
Diag(ArgLoc, diag::err_hlsl_space_on_global_constant);
2407 if (RegType == RegisterType::CBuffer)
2408 S.
Diag(ArgLoc, diag::warn_hlsl_deprecated_register_type_b);
2409 else if (RegType != RegisterType::C)
2410 S.
Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
2414 if (RegType == RegisterType::C)
2415 S.
Diag(ArgLoc, diag::warn_hlsl_register_type_c_packoffset);
2417 S.
Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
2427 S.
Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
2435 bool RegisterTypesDetected[5] = {
false};
2436 RegisterTypesDetected[
static_cast<int>(regType)] =
true;
2439 if (HLSLResourceBindingAttr *
attr =
2440 dyn_cast<HLSLResourceBindingAttr>(*it)) {
2443 if (RegisterTypesDetected[
static_cast<int>(otherRegType)]) {
2444 int otherRegTypeNum =
static_cast<int>(otherRegType);
2446 diag::err_hlsl_duplicate_register_annotation)
2450 RegisterTypesDetected[
static_cast<int>(otherRegType)] =
true;
2458 bool SpecifiedSpace) {
2463 "expecting VarDecl or HLSLBufferDecl");
2475 const uint64_t &Limit,
2478 uint64_t ArrayCount = 1) {
2483 if (StartSlot > Limit)
2487 if (
const auto *AT = dyn_cast<ArrayType>(T)) {
2490 if (
const auto *CAT = dyn_cast<ConstantArrayType>(AT))
2491 Count = CAT->
getSize().getZExtValue();
2495 ArrayCount * Count);
2499 if (
auto ResTy = dyn_cast<HLSLAttributedResourceType>(T)) {
2502 if (ResTy->getAttrs().ResourceClass != ResClass)
2506 uint64_t EndSlot = StartSlot + ArrayCount - 1;
2507 if (EndSlot > Limit)
2511 StartSlot = EndSlot + 1;
2516 if (
const auto *RT = dyn_cast<RecordType>(T)) {
2519 if (
const auto *CXXRD = dyn_cast<CXXRecordDecl>(RD)) {
2522 ResClass, Ctx, ArrayCount))
2529 ResClass, Ctx, ArrayCount))
2543 const uint64_t Limit = UINT32_MAX;
2544 if (SlotNum > Limit)
2549 if (RegTy == RegisterType::C || RegTy == RegisterType::I)
2552 if (
VarDecl *VD = dyn_cast<VarDecl>(TheDecl)) {
2553 uint64_t BaseSlot = SlotNum;
2561 return (BaseSlot > Limit);
2568 return (SlotNum > Limit);
2571 llvm_unreachable(
"unexpected decl type");
2575 if (
VarDecl *VD = dyn_cast<VarDecl>(TheDecl)) {
2577 if (
const auto *IAT = dyn_cast<IncompleteArrayType>(Ty))
2578 Ty = IAT->getElementType();
2580 diag::err_incomplete_type))
2584 StringRef Slot =
"";
2585 StringRef Space =
"";
2589 Diag(AL.
getLoc(), diag::err_attribute_argument_type)
2599 Diag(AL.
getLoc(), diag::err_attribute_argument_type)
2605 SpaceLoc = Loc->
getLoc();
2608 if (Str.starts_with(
"space")) {
2610 SpaceLoc = Loc->
getLoc();
2619 std::optional<unsigned> SlotNum;
2620 unsigned SpaceNum = 0;
2623 if (!Slot.empty()) {
2625 Diag(SlotLoc, diag::err_hlsl_binding_type_invalid) << Slot.substr(0, 1);
2628 if (RegType == RegisterType::I) {
2629 Diag(SlotLoc, diag::warn_hlsl_deprecated_register_type_i);
2632 const StringRef SlotNumStr = Slot.substr(1);
2637 if (SlotNumStr.getAsInteger(10, N)) {
2638 Diag(SlotLoc, diag::err_hlsl_unsupported_register_number);
2646 Diag(SlotLoc, diag::err_hlsl_register_number_too_large);
2655 if (!Space.starts_with(
"space")) {
2656 Diag(SpaceLoc, diag::err_hlsl_expected_space) << Space;
2659 StringRef SpaceNumStr = Space.substr(5);
2660 if (SpaceNumStr.getAsInteger(10, SpaceNum)) {
2661 Diag(SpaceLoc, diag::err_hlsl_expected_space) << Space;
2666 if (SlotNum.has_value())
2671 HLSLResourceBindingAttr *NewAttr =
2672 HLSLResourceBindingAttr::Create(
getASTContext(), Slot, Space, AL);
2674 NewAttr->setBinding(RegType, SlotNum, SpaceNum);
2729 llvm::DenseMap<const FunctionDecl *, unsigned> ScannedDecls;
2733 llvm::Triple::EnvironmentType CurrentShaderEnvironment;
2734 unsigned CurrentShaderStageBit;
2739 bool ReportOnlyShaderStageIssues;
2742 void SetShaderStageContext(llvm::Triple::EnvironmentType ShaderType) {
2743 static_assert(
sizeof(
unsigned) >= 4);
2744 assert(HLSLShaderAttr::isValidShaderType(ShaderType));
2745 assert((
unsigned)(ShaderType - llvm::Triple::Pixel) < 31 &&
2746 "ShaderType is too big for this bitmap");
2749 unsigned bitmapIndex = ShaderType - llvm::Triple::Pixel;
2750 CurrentShaderEnvironment = ShaderType;
2751 CurrentShaderStageBit = (1 << bitmapIndex);
2754 void SetUnknownShaderStageContext() {
2755 CurrentShaderEnvironment = llvm::Triple::UnknownEnvironment;
2756 CurrentShaderStageBit = (1 << 31);
2759 llvm::Triple::EnvironmentType GetCurrentShaderEnvironment()
const {
2760 return CurrentShaderEnvironment;
2763 bool InUnknownShaderStageContext()
const {
2764 return CurrentShaderEnvironment == llvm::Triple::UnknownEnvironment;
2768 void AddToScannedFunctions(
const FunctionDecl *FD) {
2769 unsigned &ScannedStages = ScannedDecls[FD];
2770 ScannedStages |= CurrentShaderStageBit;
2773 unsigned GetScannedStages(
const FunctionDecl *FD) {
return ScannedDecls[FD]; }
2775 bool WasAlreadyScannedInCurrentStage(
const FunctionDecl *FD) {
2776 return WasAlreadyScannedInCurrentStage(GetScannedStages(FD));
2779 bool WasAlreadyScannedInCurrentStage(
unsigned ScannerStages) {
2780 return ScannerStages & CurrentShaderStageBit;
2783 static bool NeverBeenScanned(
unsigned ScannedStages) {
2784 return ScannedStages == 0;
2788 void HandleFunctionOrMethodRef(FunctionDecl *FD, Expr *RefExpr);
2789 void CheckDeclAvailability(NamedDecl *D,
const AvailabilityAttr *AA,
2791 const AvailabilityAttr *FindAvailabilityAttr(
const Decl *D);
2792 bool HasMatchingEnvironmentOrNone(
const AvailabilityAttr *AA);
2795 DiagnoseHLSLAvailability(Sema &SemaRef)
2797 CurrentShaderEnvironment(llvm::Triple::UnknownEnvironment),
2798 CurrentShaderStageBit(0), ReportOnlyShaderStageIssues(
false) {}
2801 void RunOnTranslationUnit(
const TranslationUnitDecl *TU);
2802 void RunOnFunction(
const FunctionDecl *FD);
2804 bool VisitDeclRefExpr(DeclRefExpr *DRE)
override {
2805 FunctionDecl *FD = llvm::dyn_cast<FunctionDecl>(DRE->
getDecl());
2807 HandleFunctionOrMethodRef(FD, DRE);
2811 bool VisitMemberExpr(MemberExpr *ME)
override {
2812 FunctionDecl *FD = llvm::dyn_cast<FunctionDecl>(ME->
getMemberDecl());
2814 HandleFunctionOrMethodRef(FD, ME);
2819void DiagnoseHLSLAvailability::HandleFunctionOrMethodRef(
FunctionDecl *FD,
2822 "expected DeclRefExpr or MemberExpr");
2826 if (FD->
hasBody(FDWithBody)) {
2827 if (!WasAlreadyScannedInCurrentStage(FDWithBody))
2828 DeclsToScan.push_back(FDWithBody);
2833 const AvailabilityAttr *AA = FindAvailabilityAttr(FD);
2835 CheckDeclAvailability(
2839void DiagnoseHLSLAvailability::RunOnTranslationUnit(
2848 DeclContextsToScan.push_back(TU);
2850 while (!DeclContextsToScan.empty()) {
2851 const DeclContext *DC = DeclContextsToScan.pop_back_val();
2852 for (
auto &D : DC->
decls()) {
2859 if (llvm::dyn_cast<NamespaceDecl>(D) || llvm::dyn_cast<ExportDecl>(D)) {
2860 DeclContextsToScan.push_back(llvm::dyn_cast<DeclContext>(D));
2865 const FunctionDecl *FD = llvm::dyn_cast<FunctionDecl>(D);
2870 if (HLSLShaderAttr *ShaderAttr = FD->
getAttr<HLSLShaderAttr>()) {
2871 SetShaderStageContext(ShaderAttr->getType());
2880 for (
const auto *Redecl : FD->
redecls()) {
2881 if (Redecl->isInExportDeclContext()) {
2888 SetUnknownShaderStageContext();
2896void DiagnoseHLSLAvailability::RunOnFunction(
const FunctionDecl *FD) {
2897 assert(DeclsToScan.empty() &&
"DeclsToScan should be empty");
2898 DeclsToScan.push_back(FD);
2900 while (!DeclsToScan.empty()) {
2908 const unsigned ScannedStages = GetScannedStages(FD);
2909 if (WasAlreadyScannedInCurrentStage(ScannedStages))
2912 ReportOnlyShaderStageIssues = !NeverBeenScanned(ScannedStages);
2914 AddToScannedFunctions(FD);
2919bool DiagnoseHLSLAvailability::HasMatchingEnvironmentOrNone(
2920 const AvailabilityAttr *AA) {
2925 llvm::Triple::EnvironmentType CurrentEnv = GetCurrentShaderEnvironment();
2926 if (CurrentEnv == llvm::Triple::UnknownEnvironment)
2929 llvm::Triple::EnvironmentType AttrEnv =
2930 AvailabilityAttr::getEnvironmentType(IIEnvironment->
getName());
2932 return CurrentEnv == AttrEnv;
2935const AvailabilityAttr *
2936DiagnoseHLSLAvailability::FindAvailabilityAttr(
const Decl *D) {
2937 AvailabilityAttr
const *PartialMatch =
nullptr;
2941 for (
const auto *A : D->
attrs()) {
2942 if (
const auto *Avail = dyn_cast<AvailabilityAttr>(A)) {
2943 const AvailabilityAttr *EffectiveAvail = Avail->getEffectiveAttr();
2944 StringRef AttrPlatform = EffectiveAvail->getPlatform()->getName();
2945 StringRef TargetPlatform =
2949 if (AttrPlatform == TargetPlatform) {
2951 if (HasMatchingEnvironmentOrNone(EffectiveAvail))
2953 PartialMatch = Avail;
2957 return PartialMatch;
2962void DiagnoseHLSLAvailability::CheckDeclAvailability(
NamedDecl *D,
2963 const AvailabilityAttr *AA,
2982 if (ReportOnlyShaderStageIssues)
2988 if (InUnknownShaderStageContext())
2993 bool EnvironmentMatches = HasMatchingEnvironmentOrNone(AA);
2994 VersionTuple Introduced = AA->getIntroduced();
3003 llvm::StringRef PlatformName(
3006 llvm::StringRef CurrentEnvStr =
3007 llvm::Triple::getEnvironmentTypeName(GetCurrentShaderEnvironment());
3009 llvm::StringRef AttrEnvStr =
3010 AA->getEnvironment() ? AA->getEnvironment()->getName() :
"";
3011 bool UseEnvironment = !AttrEnvStr.empty();
3013 if (EnvironmentMatches) {
3014 SemaRef.
Diag(
Range.getBegin(), diag::warn_hlsl_availability)
3015 <<
Range << D << PlatformName << Introduced.getAsString()
3016 << UseEnvironment << CurrentEnvStr;
3018 SemaRef.
Diag(
Range.getBegin(), diag::warn_hlsl_availability_unavailable)
3022 SemaRef.
Diag(D->
getLocation(), diag::note_partial_availability_specified_here)
3023 << D << PlatformName << Introduced.getAsString()
3025 << UseEnvironment << AttrEnvStr << CurrentEnvStr;
3032 if (!DefaultCBufferDecls.empty()) {
3035 DefaultCBufferDecls);
3038 SemaRef.getCurLexicalContext()->addDecl(DefaultCBuffer);
3042 for (
const Decl *VD : DefaultCBufferDecls) {
3043 const HLSLResourceBindingAttr *RBA =
3044 VD->
getAttr<HLSLResourceBindingAttr>();
3045 if (RBA && RBA->hasRegisterSlot() &&
3046 RBA->getRegisterType() == HLSLResourceBindingAttr::RegisterType::C) {
3053 SemaRef.Consumer.HandleTopLevelDecl(DG);
3055 diagnoseAvailabilityViolations(TU);
3064 "expected member expr to have resource record type or array of them");
3070 const Expr *NonConstIndexExpr =
nullptr;
3073 if (
const DeclRefExpr *DRE = dyn_cast<DeclRefExpr>(E)) {
3074 if (!NonConstIndexExpr)
3082 diag::err_hlsl_resource_member_array_access_not_constant);
3086 if (
const auto *ASE = dyn_cast<ArraySubscriptExpr>(E)) {
3087 const Expr *IdxExpr = ASE->getIdx();
3089 NonConstIndexExpr = IdxExpr;
3091 }
else if (
const auto *SubME = dyn_cast<MemberExpr>(E)) {
3092 E = SubME->getBase();
3093 }
else if (
const auto *ICE = dyn_cast<ImplicitCastExpr>(E)) {
3094 E = ICE->getSubExpr();
3096 llvm_unreachable(
"unexpected expr type in resource member access");
3109 TI.
getTriple().getEnvironment() != llvm::Triple::EnvironmentType::Library)
3112 DiagnoseHLSLAvailability(
SemaRef).RunOnTranslationUnit(TU);
3119 for (
unsigned I = 1, N = TheCall->
getNumArgs(); I < N; ++I) {
3122 S->
Diag(TheCall->
getBeginLoc(), diag::err_vec_builtin_incompatible_vector)
3147 for (
unsigned I = 0; I < TheCall->
getNumArgs(); ++I) {
3162 if (!BaseType->isFloat32Type())
3163 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3164 << ArgOrdinal << 5 << 0
3176 if (!BaseType->isHalfType() && !BaseType->isFloat32Type())
3177 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3178 << ArgOrdinal << 5 << 0
3192 if (!BaseType->isDoubleType()) {
3195 return S->
Diag(Loc, diag::err_builtin_requires_double_type)
3196 << ArgOrdinal << PassedType;
3203 unsigned ArgIndex) {
3204 auto *Arg = TheCall->
getArg(ArgIndex);
3206 if (Arg->IgnoreCasts()->isModifiableLvalue(S->
Context, &OrigLoc) ==
3209 S->
Diag(OrigLoc, diag::error_hlsl_inout_lvalue) << Arg << 0;
3219 if (VecTy->getElementType()->isDoubleType())
3220 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3221 << ArgOrdinal << 1 << 0 << 1
3231 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3232 << ArgOrdinal << 5 << 1
3241 if (VecTy->getElementType()->isUnsignedIntegerType())
3244 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3245 << ArgOrdinal << 4 << 3 << 0
3254 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3255 << ArgOrdinal << 5 << 3
3261 unsigned ArgOrdinal,
unsigned Width) {
3264 ArgTy = VTy->getElementType();
3266 uint64_t ElementBitCount =
3268 if (ElementBitCount != Width) {
3270 diag::err_integer_incorrect_bit_count)
3271 << Width << ElementBitCount;
3288 unsigned ArgIndex) {
3297 diag::err_typecheck_expect_scalar_or_vector)
3298 << ArgType << Scalar;
3305 QualType Scalar,
unsigned ArgIndex) {
3316 if (
const auto *VTy = ArgType->getAs<
VectorType>()) {
3329 diag::err_typecheck_expect_scalar_or_vector_or_matrix)
3330 << ArgType << Scalar;
3335 unsigned ArgIndex) {
3340 if (!(ArgType->isScalarType() ||
3341 (VTy && VTy->getElementType()->isScalarType()))) {
3343 diag::err_typecheck_expect_any_scalar_or_vector)
3353 unsigned ArgIndex) {
3355 assert(ArgIndex < TheCall->getNumArgs());
3363 diag::err_typecheck_expect_any_scalar_or_vector)
3388 diag::err_typecheck_call_different_arg_types)
3407 Arg1ScalarTy = VTy->getElementType();
3411 Arg2ScalarTy = VTy->getElementType();
3414 S->
Diag(Arg1->
getBeginLoc(), diag::err_hlsl_builtin_scalar_vector_mismatch)
3415 << 1 << TheCall->
getCallee() << Arg1Ty << Arg2Ty;
3425 if (Arg1Length > 0 && Arg0Length != Arg1Length) {
3427 diag::err_typecheck_vector_lengths_not_equal)
3433 if (Arg2Length > 0 && Arg0Length != Arg2Length) {
3435 diag::err_typecheck_vector_lengths_not_equal)
3447 assert(TheCall->
getNumArgs() > IndexArgIndex &&
"Index argument missing");
3450 unsigned int ActualDim = 1;
3452 ActualDim = VTy->getNumElements();
3453 IndexTy = VTy->getElementType();
3457 diag::err_typecheck_expect_int)
3463 const HLSLAttributedResourceType *ResTy =
3465 assert(ResTy &&
"Resource argument must be a resource");
3466 HLSLAttributedResourceType::Attributes ResAttrs = ResTy->getAttrs();
3468 unsigned int ExpectedDim = 1;
3469 if (ResAttrs.ResourceDimension != llvm::dxil::ResourceDimension::Unknown)
3472 if (ActualDim != ExpectedDim) {
3474 diag::err_hlsl_builtin_resource_coordinate_dimension_mismatch)
3485 llvm::function_ref<
bool(
const HLSLAttributedResourceType *ResType)> Check =
3489 const HLSLAttributedResourceType *ResTy =
3493 diag::err_typecheck_expect_hlsl_resource)
3497 if (Check && Check(ResTy)) {
3499 diag::err_invalid_hlsl_resource_type)
3507 QualType BaseType,
unsigned ExpectedCount,
3509 unsigned PassedCount = 1;
3511 PassedCount = VecTy->getNumElements();
3513 if (PassedCount != ExpectedCount) {
3516 S->
Diag(Loc, diag::err_typecheck_convert_incompatible)
3528 [](
const HLSLAttributedResourceType *ResType) {
3529 return ResType->getAttrs().ResourceDimension ==
3530 llvm::dxil::ResourceDimension::Unknown;
3536 [](
const HLSLAttributedResourceType *ResType) {
3537 return ResType->getAttrs().ResourceClass !=
3538 llvm::hlsl::ResourceClass::Sampler;
3546 unsigned ExpectedDim =
3574 unsigned NextIdx = 3;
3580 diag::err_typecheck_convert_incompatible)
3588 Expr *ComponentArg = TheCall->
getArg(NextIdx);
3592 diag::err_typecheck_convert_incompatible)
3599 std::optional<llvm::APSInt> ComponentOpt =
3602 int64_t ComponentVal = ComponentOpt->getSExtValue();
3603 if (ComponentVal != 0) {
3606 assert(ComponentVal >= 0 && ComponentVal <= 3 &&
3607 "The component is not in the expected range.");
3609 diag::err_hlsl_gathercmp_invalid_component)
3619 const HLSLAttributedResourceType *ResourceTy =
3622 unsigned ExpectedDim =
3631 assert(ResourceTy->hasContainedType() &&
3632 "Expecting a contained type for resource with a dimension "
3634 QualType ReturnType = ResourceTy->getContainedType();
3638 S.
Diag(TheCall->
getBeginLoc(), diag::err_hlsl_samplecmp_requires_float);
3644 ReturnType = VecTy->getElementType();
3657 [](
const HLSLAttributedResourceType *ResType) {
3658 return ResType->getAttrs().ResourceDimension ==
3659 llvm::dxil::ResourceDimension::Unknown;
3667 unsigned ExpectedDim =
3676 EltTy = VTy->getElementType();
3691 TheCall->
setType(ResourceTy->getContainedType());
3696 unsigned MinArgs, MaxArgs;
3724 const HLSLAttributedResourceType *ResourceTy =
3726 unsigned ExpectedDim =
3729 unsigned NextIdx = 3;
3738 diag::err_typecheck_convert_incompatible)
3773 diag::err_typecheck_convert_incompatible)
3779 assert(ResourceTy->hasContainedType() &&
3780 "Expecting a contained type for resource with a dimension "
3782 QualType ReturnType = ResourceTy->getContainedType();
3785 S.
Diag(TheCall->
getBeginLoc(), diag::err_hlsl_samplecmp_requires_float);
3798 switch (BuiltinID) {
3799 case Builtin::BI__builtin_hlsl_adduint64: {
3800 if (
SemaRef.checkArgCount(TheCall, 2))
3814 if (NumElementsArg != 2 && NumElementsArg != 4) {
3816 << 1 << 64 << NumElementsArg * 32;
3830 case Builtin::BI__builtin_hlsl_resource_getpointer: {
3831 if (
SemaRef.checkArgCount(TheCall, 2) ||
3838 QualType ContainedTy = ResourceTy->getContainedType();
3841 ReturnType =
SemaRef.Context.getPointerType(ReturnType);
3847 case Builtin::BI__builtin_hlsl_resource_getpointer_typed: {
3848 if (
SemaRef.checkArgCount(TheCall, 3) ||
3855 "expected pointer type for second argument");
3862 diag::err_invalid_use_of_array_type);
3866 ReturnType =
SemaRef.Context.getPointerType(ReturnType);
3871 case Builtin::BI__builtin_hlsl_resource_load_with_status: {
3872 if (
SemaRef.checkArgCount(TheCall, 3) ||
3875 SemaRef.getASTContext().UnsignedIntTy) ||
3877 SemaRef.getASTContext().UnsignedIntTy) ||
3883 QualType ReturnType = ResourceTy->getContainedType();
3888 case Builtin::BI__builtin_hlsl_resource_load_with_status_typed: {
3889 if (
SemaRef.checkArgCount(TheCall, 4) ||
3892 SemaRef.getASTContext().UnsignedIntTy) ||
3894 SemaRef.getASTContext().UnsignedIntTy) ||
3900 "expected pointer type for second argument");
3907 diag::err_invalid_use_of_array_type);
3913 case Builtin::BI__builtin_hlsl_resource_load_level:
3915 case Builtin::BI__builtin_hlsl_resource_sample:
3917 case Builtin::BI__builtin_hlsl_resource_sample_bias:
3919 case Builtin::BI__builtin_hlsl_resource_sample_grad:
3921 case Builtin::BI__builtin_hlsl_resource_sample_level:
3923 case Builtin::BI__builtin_hlsl_resource_sample_cmp:
3925 case Builtin::BI__builtin_hlsl_resource_sample_cmp_level_zero:
3927 case Builtin::BI__builtin_hlsl_resource_calculate_lod:
3928 case Builtin::BI__builtin_hlsl_resource_calculate_lod_unclamped:
3930 case Builtin::BI__builtin_hlsl_resource_gather:
3932 case Builtin::BI__builtin_hlsl_resource_gather_cmp:
3934 case Builtin::BI__builtin_hlsl_resource_uninitializedhandle: {
3935 assert(TheCall->
getNumArgs() == 1 &&
"expected 1 arg");
3941 case Builtin::BI__builtin_hlsl_resource_handlefrombinding: {
3942 assert(TheCall->
getNumArgs() == 6 &&
"expected 6 args");
3948 case Builtin::BI__builtin_hlsl_resource_handlefromimplicitbinding: {
3949 assert(TheCall->
getNumArgs() == 6 &&
"expected 6 args");
3955 case Builtin::BI__builtin_hlsl_resource_counterhandlefromimplicitbinding: {
3956 assert(TheCall->
getNumArgs() == 3 &&
"expected 3 args");
3959 auto *MainResType = MainHandleTy->
getAs<HLSLAttributedResourceType>();
3960 auto MainAttrs = MainResType->getAttrs();
3961 assert(!MainAttrs.IsCounter &&
"cannot create a counter from a counter");
3962 MainAttrs.IsCounter =
true;
3964 MainResType->getWrappedType(), MainResType->getContainedType(),
3968 TheCall->
setType(CounterHandleTy);
3971 case Builtin::BI__builtin_hlsl_and:
3972 case Builtin::BI__builtin_hlsl_or: {
3973 if (
SemaRef.checkArgCount(TheCall, 2))
3987 case Builtin::BI__builtin_hlsl_all:
3988 case Builtin::BI__builtin_hlsl_any: {
3989 if (
SemaRef.checkArgCount(TheCall, 1))
3995 case Builtin::BI__builtin_hlsl_asdouble: {
3996 if (
SemaRef.checkArgCount(TheCall, 2))
4000 SemaRef.Context.UnsignedIntTy,
4005 SemaRef.Context.UnsignedIntTy,
4014 case Builtin::BI__builtin_hlsl_elementwise_clamp: {
4015 if (
SemaRef.BuiltinElementwiseTernaryMath(
4021 case Builtin::BI__builtin_hlsl_dot: {
4023 if (
SemaRef.BuiltinVectorToScalarMath(TheCall))
4029 case Builtin::BI__builtin_hlsl_elementwise_firstbithigh:
4030 case Builtin::BI__builtin_hlsl_elementwise_firstbitlow: {
4031 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
4041 EltTy = VecTy->getElementType();
4042 ResTy =
SemaRef.Context.getExtVectorType(ResTy, VecTy->getNumElements());
4055 case Builtin::BI__builtin_hlsl_select: {
4056 if (
SemaRef.checkArgCount(TheCall, 3))
4064 if (VTy && VTy->getElementType()->isBooleanType() &&
4069 case Builtin::BI__builtin_hlsl_elementwise_saturate:
4070 case Builtin::BI__builtin_hlsl_elementwise_rcp: {
4071 if (
SemaRef.checkArgCount(TheCall, 1))
4077 diag::err_builtin_invalid_arg_type)
4080 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
4084 case Builtin::BI__builtin_hlsl_elementwise_degrees:
4085 case Builtin::BI__builtin_hlsl_elementwise_radians:
4086 case Builtin::BI__builtin_hlsl_elementwise_rsqrt:
4087 case Builtin::BI__builtin_hlsl_elementwise_frac:
4088 case Builtin::BI__builtin_hlsl_elementwise_ddx_coarse:
4089 case Builtin::BI__builtin_hlsl_elementwise_ddy_coarse:
4090 case Builtin::BI__builtin_hlsl_elementwise_ddx_fine:
4091 case Builtin::BI__builtin_hlsl_elementwise_ddy_fine: {
4092 if (
SemaRef.checkArgCount(TheCall, 1))
4097 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
4101 case Builtin::BI__builtin_hlsl_elementwise_isinf:
4102 case Builtin::BI__builtin_hlsl_elementwise_isnan: {
4103 if (
SemaRef.checkArgCount(TheCall, 1))
4108 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
4113 case Builtin::BI__builtin_hlsl_lerp: {
4114 if (
SemaRef.checkArgCount(TheCall, 3))
4121 if (
SemaRef.BuiltinElementwiseTernaryMath(TheCall))
4125 case Builtin::BI__builtin_hlsl_mad: {
4126 if (
SemaRef.BuiltinElementwiseTernaryMath(
4132 case Builtin::BI__builtin_hlsl_mul: {
4133 if (
SemaRef.checkArgCount(TheCall, 2))
4142 if (
const auto *VTy = T->getAs<
VectorType>())
4143 return VTy->getElementType();
4145 return MTy->getElementType();
4149 QualType EltTy0 = getElemType(Ty0);
4158 if (IsVec0 && IsMat1) {
4161 }
else if (IsMat0 && IsVec1) {
4165 assert(IsMat0 && IsMat1);
4175 case Builtin::BI__builtin_hlsl_normalize: {
4176 if (
SemaRef.checkArgCount(TheCall, 1))
4187 case Builtin::BI__builtin_elementwise_fma: {
4188 if (
SemaRef.checkArgCount(TheCall, 3) ||
4203 case Builtin::BI__builtin_hlsl_transpose: {
4204 if (
SemaRef.checkArgCount(TheCall, 1))
4213 << 1 << 3 << 0 << 0 << ArgTy;
4218 MatTy->getElementType(), MatTy->getNumColumns(), MatTy->getNumRows());
4222 case Builtin::BI__builtin_hlsl_elementwise_sign: {
4223 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
4231 case Builtin::BI__builtin_hlsl_step: {
4232 if (
SemaRef.checkArgCount(TheCall, 2))
4244 case Builtin::BI__builtin_hlsl_wave_active_all_equal: {
4245 if (
SemaRef.checkArgCount(TheCall, 1))
4259 unsigned NumElts = VecTy->getNumElements();
4269 case Builtin::BI__builtin_hlsl_wave_active_max:
4270 case Builtin::BI__builtin_hlsl_wave_active_min:
4271 case Builtin::BI__builtin_hlsl_wave_active_sum:
4272 case Builtin::BI__builtin_hlsl_wave_active_product: {
4273 if (
SemaRef.checkArgCount(TheCall, 1))
4286 case Builtin::BI__builtin_hlsl_wave_active_bit_or:
4287 case Builtin::BI__builtin_hlsl_wave_active_bit_xor:
4288 case Builtin::BI__builtin_hlsl_wave_active_bit_and: {
4289 if (
SemaRef.checkArgCount(TheCall, 1))
4304 (VTy && VTy->getElementType()->isIntegerType()))) {
4306 diag::err_builtin_invalid_arg_type)
4307 << ArgTyExpr <<
SemaRef.Context.UnsignedIntTy << 1 << 0 << 0;
4317 case Builtin::BI__builtin_elementwise_bitreverse: {
4325 case Builtin::BI__builtin_hlsl_wave_prefix_count_bits: {
4326 if (
SemaRef.checkArgCount(TheCall, 1))
4331 if (!(
ArgType->isScalarType())) {
4333 diag::err_typecheck_expect_any_scalar_or_vector)
4338 if (!(
ArgType->isBooleanType())) {
4340 diag::err_typecheck_expect_any_scalar_or_vector)
4347 case Builtin::BI__builtin_hlsl_wave_read_lane_at: {
4348 if (
SemaRef.checkArgCount(TheCall, 2))
4356 diag::err_typecheck_convert_incompatible)
4357 << ArgTyIndex <<
SemaRef.Context.UnsignedIntTy << 1 << 0 << 0;
4370 case Builtin::BI__builtin_hlsl_wave_get_lane_index: {
4371 if (
SemaRef.checkArgCount(TheCall, 0))
4375 case Builtin::BI__builtin_hlsl_wave_prefix_sum:
4376 case Builtin::BI__builtin_hlsl_wave_prefix_product: {
4377 if (
SemaRef.checkArgCount(TheCall, 1))
4390 case Builtin::BI__builtin_hlsl_quad_read_across_x:
4391 case Builtin::BI__builtin_hlsl_quad_read_across_y: {
4392 if (
SemaRef.checkArgCount(TheCall, 1))
4404 case Builtin::BI__builtin_hlsl_elementwise_splitdouble: {
4405 if (
SemaRef.checkArgCount(TheCall, 3))
4420 case Builtin::BI__builtin_hlsl_elementwise_clip: {
4421 if (
SemaRef.checkArgCount(TheCall, 1))
4428 case Builtin::BI__builtin_elementwise_acos:
4429 case Builtin::BI__builtin_elementwise_asin:
4430 case Builtin::BI__builtin_elementwise_atan:
4431 case Builtin::BI__builtin_elementwise_atan2:
4432 case Builtin::BI__builtin_elementwise_ceil:
4433 case Builtin::BI__builtin_elementwise_cos:
4434 case Builtin::BI__builtin_elementwise_cosh:
4435 case Builtin::BI__builtin_elementwise_exp:
4436 case Builtin::BI__builtin_elementwise_exp2:
4437 case Builtin::BI__builtin_elementwise_exp10:
4438 case Builtin::BI__builtin_elementwise_floor:
4439 case Builtin::BI__builtin_elementwise_fmod:
4440 case Builtin::BI__builtin_elementwise_log:
4441 case Builtin::BI__builtin_elementwise_log2:
4442 case Builtin::BI__builtin_elementwise_log10:
4443 case Builtin::BI__builtin_elementwise_pow:
4444 case Builtin::BI__builtin_elementwise_roundeven:
4445 case Builtin::BI__builtin_elementwise_sin:
4446 case Builtin::BI__builtin_elementwise_sinh:
4447 case Builtin::BI__builtin_elementwise_sqrt:
4448 case Builtin::BI__builtin_elementwise_tan:
4449 case Builtin::BI__builtin_elementwise_tanh:
4450 case Builtin::BI__builtin_elementwise_trunc: {
4456 case Builtin::BI__builtin_hlsl_buffer_update_counter: {
4457 assert(TheCall->
getNumArgs() == 2 &&
"expected 2 args");
4458 auto checkResTy = [](
const HLSLAttributedResourceType *ResTy) ->
bool {
4459 return !(ResTy->getAttrs().ResourceClass == ResourceClass::UAV &&
4460 ResTy->getAttrs().RawBuffer && ResTy->hasContainedType());
4465 std::optional<llvm::APSInt> Offset =
4467 if (!Offset.has_value() ||
std::abs(Offset->getExtValue()) != 1) {
4469 diag::err_hlsl_expect_arg_const_int_one_or_neg_one)
4475 case Builtin::BI__builtin_hlsl_elementwise_f16tof32: {
4476 if (
SemaRef.checkArgCount(TheCall, 1))
4487 ArgTy = VTy->getElementType();
4490 diag::err_builtin_invalid_arg_type)
4499 case Builtin::BI__builtin_hlsl_elementwise_f32tof16: {
4500 if (
SemaRef.checkArgCount(TheCall, 1))
4515 WorkList.push_back(BaseTy);
4516 while (!WorkList.empty()) {
4517 QualType T = WorkList.pop_back_val();
4518 T = T.getCanonicalType().getUnqualifiedType();
4519 if (
const auto *AT = dyn_cast<ConstantArrayType>(T)) {
4527 for (uint64_t Ct = 0; Ct < AT->
getZExtSize(); ++Ct)
4528 llvm::append_range(List, ElementFields);
4533 if (
const auto *VT = dyn_cast<VectorType>(T)) {
4534 List.insert(List.end(), VT->getNumElements(), VT->getElementType());
4537 if (
const auto *MT = dyn_cast<ConstantMatrixType>(T)) {
4538 List.insert(List.end(), MT->getNumElementsFlattened(),
4539 MT->getElementType());
4542 if (
const auto *RD = T->getAsCXXRecordDecl()) {
4543 if (RD->isStandardLayout())
4544 RD = RD->getStandardLayoutBaseWithFields();
4548 if (RD->
isUnion() || !RD->isAggregate()) {
4554 for (
const auto *FD : RD->
fields())
4555 if (!FD->isUnnamedBitField())
4556 FieldTypes.push_back(FD->
getType());
4558 std::reverse(FieldTypes.begin(), FieldTypes.end());
4559 llvm::append_range(WorkList, FieldTypes);
4563 if (!RD->isStandardLayout()) {
4565 for (
const auto &
Base : RD->bases())
4566 FieldTypes.push_back(
Base.getType());
4567 std::reverse(FieldTypes.begin(), FieldTypes.end());
4568 llvm::append_range(WorkList, FieldTypes);
4590 if (
SemaRef.Context.getTypeSize(QT) / 8 > 16)
4596 int ArraySize = VT->getNumElements();
4601 QualType ElTy = VT->getElementType();
4605 if (
SemaRef.Context.getTypeSize(QT) / 8 > 16)
4621 if (
SemaRef.getASTContext().hasSameType(T1, T2))
4630 return llvm::equal(T1Types, T2Types,
4632 return SemaRef.IsLayoutCompatible(LHS, RHS);
4641 bool HadError =
false;
4643 for (
unsigned i = 0, e =
New->getNumParams(); i != e; ++i) {
4651 const auto *NDAttr = NewParam->
getAttr<HLSLParamModifierAttr>();
4652 unsigned NSpellingIdx = (NDAttr ? NDAttr->getSpellingListIndex() : 0);
4653 const auto *ODAttr = OldParam->
getAttr<HLSLParamModifierAttr>();
4654 unsigned OSpellingIdx = (ODAttr ? ODAttr->getSpellingListIndex() : 0);
4656 if (NSpellingIdx != OSpellingIdx) {
4658 diag::err_hlsl_param_qualifier_mismatch)
4659 << NDAttr << NewParam;
4675 if (
SemaRef.getASTContext().hasSameUnqualifiedType(SrcTy, DestTy))
4690 llvm_unreachable(
"HLSL doesn't support pointers.");
4693 llvm_unreachable(
"HLSL doesn't support complex types.");
4695 llvm_unreachable(
"HLSL doesn't support fixed point types.");
4697 llvm_unreachable(
"Should have returned before this");
4707 llvm_unreachable(
"HLSL doesn't support complex types.");
4709 llvm_unreachable(
"HLSL doesn't support fixed point types.");
4714 llvm_unreachable(
"HLSL doesn't support pointers.");
4716 llvm_unreachable(
"Should have returned before this");
4722 llvm_unreachable(
"HLSL doesn't support pointers.");
4725 llvm_unreachable(
"HLSL doesn't support fixed point types.");
4729 llvm_unreachable(
"HLSL doesn't support complex types.");
4732 llvm_unreachable(
"Unhandled scalar cast");
4753 !(SrcMatTy && SrcMatTy->getNumElementsFlattened() == 1))
4759 SrcTy = SrcMatTy->getElementType();
4764 for (
unsigned I = 0, Size = DestTypes.size(); I < Size; ++I) {
4765 if (DestTypes[I]->isUnionType())
4797 if (SrcTypes.size() < DestTypes.size())
4800 unsigned SrcSize = SrcTypes.size();
4801 unsigned DstSize = DestTypes.size();
4803 for (I = 0; I < DstSize && I < SrcSize; I++) {
4804 if (SrcTypes[I]->isUnionType() || DestTypes[I]->isUnionType())
4812 for (; I < SrcSize; I++) {
4813 if (SrcTypes[I]->isUnionType())
4820 assert(Param->hasAttr<HLSLParamModifierAttr>() &&
4821 "We should not get here without a parameter modifier expression");
4822 const auto *
Attr = Param->getAttr<HLSLParamModifierAttr>();
4829 << Arg << (IsInOut ? 1 : 0);
4835 QualType Ty = Param->getType().getNonLValueExprType(Ctx);
4842 << Arg << (IsInOut ? 1 : 0);
4854 SemaRef.PerformCopyInitialization(Entity, Param->getBeginLoc(), ArgOpV);
4860 auto *OpV =
new (Ctx)
4865 Res =
SemaRef.ActOnBinOp(
SemaRef.getCurScope(), Param->getBeginLoc(),
4866 tok::equal, ArgOpV, OpV);
4882 "Pointer and reference types cannot be inout or out parameters");
4883 Ty =
SemaRef.getASTContext().getLValueReferenceType(Ty);
4899 for (
const auto *FD : RD->
fields()) {
4903 assert(RD->getNumBases() <= 1 &&
4904 "HLSL doesn't support multiple inheritance");
4905 return RD->getNumBases()
4910 if (
const auto *AT = dyn_cast<ArrayType>(Ty)) {
4911 if (
const auto *CAT = dyn_cast<ConstantArrayType>(AT))
4923 bool IsVKPushConstant = IsVulkan && VD->
hasAttr<HLSLVkPushConstantAttr>();
4928 !VD->
hasAttr<HLSLVkConstantIdAttr>() && !IsVKPushConstant &&
4934 if (
Decl->getType().hasAddressSpace())
4937 if (
Decl->getType()->isDependentType())
4949 if (
Decl->
hasAttr<HLSLVkExtBuiltinOutputAttr>()) {
4963 llvm::Triple::Vulkan;
4964 if (IsVulkan &&
Decl->
hasAttr<HLSLVkPushConstantAttr>()) {
4965 if (HasDeclaredAPushConstant)
4971 HasDeclaredAPushConstant =
true;
4998class StructBindingContext {
5001 HLSLResourceBindingAttr *RegBindingsAttrs[4];
5002 unsigned RegBindingOffset[4];
5005 static_assert(
static_cast<unsigned>(RegisterType::SRV) == 0 &&
5006 static_cast<unsigned>(RegisterType::UAV) == 1 &&
5007 static_cast<unsigned>(RegisterType::CBuffer) == 2 &&
5008 static_cast<unsigned>(RegisterType::Sampler) == 3,
5009 "unexpected register type values");
5012 HLSLVkBindingAttr *VkBindingAttr;
5013 unsigned VkBindingOffset;
5018 StructBindingContext(
VarDecl *VD) {
5019 for (
unsigned i = 0; i < 4; ++i) {
5020 RegBindingsAttrs[i] =
nullptr;
5021 RegBindingOffset[i] = 0;
5023 VkBindingAttr =
nullptr;
5024 VkBindingOffset = 0;
5030 if (
auto *RBA = dyn_cast<HLSLResourceBindingAttr>(A)) {
5032 unsigned RegTypeIdx =
static_cast<unsigned>(RegType);
5035 RegBindingsAttrs[RegTypeIdx] = RBA;
5040 if (
auto *VBA = dyn_cast<HLSLVkBindingAttr>(A))
5041 VkBindingAttr = VBA;
5048 Attr *createBindingAttr(SemaHLSL &S, ASTContext &AST,
RegisterType RegType,
5049 unsigned Range,
bool HasCounter) {
5050 assert(
static_cast<unsigned>(RegType) < 4 &&
"unexpected register type");
5052 if (VkBindingAttr) {
5053 unsigned Offset = VkBindingOffset;
5054 VkBindingOffset +=
Range;
5055 return HLSLVkBindingAttr::CreateImplicit(
5056 AST, VkBindingAttr->getBinding() + Offset, VkBindingAttr->getSet(),
5057 VkBindingAttr->getRange());
5060 HLSLResourceBindingAttr *RBA =
5061 RegBindingsAttrs[
static_cast<unsigned>(RegType)];
5062 HLSLResourceBindingAttr *NewAttr =
nullptr;
5064 if (RBA && RBA->hasRegisterSlot()) {
5067 unsigned Offset = RegBindingOffset[
static_cast<unsigned>(RegType)];
5068 RegBindingOffset[
static_cast<unsigned>(RegType)] += Range;
5070 unsigned NewSlotNumber = RBA->getSlotNumber() + Offset;
5071 StringRef NewSlotNumberStr =
5073 NewAttr = HLSLResourceBindingAttr::CreateImplicit(
5074 AST, NewSlotNumberStr, RBA->getSpace(), RBA->getRange());
5075 NewAttr->setBinding(RegType, NewSlotNumber, RBA->getSpaceNumber());
5079 NewAttr = HLSLResourceBindingAttr::CreateImplicit(AST,
"",
"0", {});
5080 NewAttr->setBinding(RegType, std::nullopt,
5081 RBA ? RBA->getSpaceNumber() : 0);
5085 NewAttr->setImplicitCounterBindingOrderID(
5094static void createGlobalResourceDeclForStruct(
5096 QualType ResTy, StructBindingContext &BindingCtx) {
5098 "expected resource type or array of resources");
5109 while (
const auto *AT = dyn_cast<ArrayType>(SingleResTy)) {
5110 const auto *CAT = dyn_cast<ConstantArrayType>(AT);
5115 const HLSLAttributedResourceType *ResHandleTy =
5116 HLSLAttributedResourceType::findHandleTypeOnResource(SingleResTy);
5120 Attr *BindingAttr = BindingCtx.createBindingAttr(
5122 ResDecl->
addAttr(BindingAttr);
5123 ResDecl->
addAttr(InternalLinkageAttr::CreateImplicit(AST));
5132 HLSLAssociatedResourceDeclAttr::CreateImplicit(AST, ResDecl));
5139static void handleArrayOfStructWithResources(
5141 EmbeddedResourceNameBuilder &NameBuilder, StructBindingContext &BindingCtx);
5146static void handleStructWithResources(
Sema &S,
VarDecl *ParentVD,
5148 EmbeddedResourceNameBuilder &NameBuilder,
5149 StructBindingContext &BindingCtx) {
5152 assert(RD->
getNumBases() <= 1 &&
"HLSL doesn't support multiple inheritance");
5159 handleStructWithResources(S, ParentVD, BaseRD, NameBuilder, BindingCtx);
5173 createGlobalResourceDeclForStruct(S, ParentVD, FD->
getLocation(), II,
5176 handleStructWithResources(S, ParentVD, RD, NameBuilder, BindingCtx);
5178 }
else if (
const auto *ArrayTy = dyn_cast<ConstantArrayType>(FDTy)) {
5180 "resource arrays should have been already handled");
5181 handleArrayOfStructWithResources(S, ParentVD, ArrayTy, NameBuilder,
5190handleArrayOfStructWithResources(
Sema &S,
VarDecl *ParentVD,
5192 EmbeddedResourceNameBuilder &NameBuilder,
5193 StructBindingContext &BindingCtx) {
5201 if (!SubCAT && !ElementRD)
5204 for (
unsigned I = 0, E = CAT->
getSize().getZExtValue(); I < E; ++I) {
5207 handleStructWithResources(S, ParentVD, ElementRD, NameBuilder,
5210 handleArrayOfStructWithResources(S, ParentVD, SubCAT, NameBuilder,
5223void SemaHLSL::handleGlobalStructOrArrayOfWithResources(
VarDecl *VD) {
5224 EmbeddedResourceNameBuilder NameBuilder(VD->
getName());
5225 StructBindingContext BindingCtx(VD);
5229 "Expected non-resource struct or array type");
5232 handleStructWithResources(
SemaRef, VD, RD, NameBuilder, BindingCtx);
5236 if (
const auto *CAT = dyn_cast<ConstantArrayType>(VDTy)) {
5237 handleArrayOfStructWithResources(
SemaRef, VD, CAT, NameBuilder, BindingCtx);
5245 if (
SemaRef.RequireCompleteType(
5248 diag::err_typecheck_decl_incomplete_type)) {
5262 DefaultCBufferDecls.push_back(VD);
5267 collectResourceBindingsOnVarDecl(VD);
5269 if (VD->
hasAttr<HLSLVkConstantIdAttr>())
5281 processExplicitBindingsOnDecl(VD);
5319 handleGlobalStructOrArrayOfWithResources(VD);
5323 if (VD->
hasAttr<HLSLGroupSharedAddressSpaceAttr>())
5332 "expected resource record type");
5348 const char *CreateMethodName;
5350 CreateMethodName = HasCounter ?
"__createFromBindingWithImplicitCounter"
5351 :
"__createFromBinding";
5353 CreateMethodName = HasCounter
5354 ?
"__createFromImplicitBindingWithImplicitCounter"
5355 :
"__createFromImplicitBinding";
5360 if (!CreateMethod) {
5365 "create method lookup should always succeed for built-in resource "
5374 Args.push_back(RegSlot);
5382 Args.push_back(OrderId);
5388 Args.push_back(Space);
5392 Args.push_back(RangeSize);
5396 Args.push_back(Index);
5398 StringRef VarName = VD->
getName();
5406 Args.push_back(NameCast);
5414 Args.push_back(CounterId);
5437 SemaRef.CheckCompleteVariableDeclaration(VD);
5443 "expected array of resource records");
5464 lookupMethod(
SemaRef, ResourceDecl,
5465 HasCounter ?
"__createFromBindingWithImplicitCounter"
5466 :
"__createFromBinding",
5470 CreateMethod = lookupMethod(
5472 HasCounter ?
"__createFromImplicitBindingWithImplicitCounter"
5473 :
"__createFromImplicitBinding",
5505std::optional<const DeclBindingInfo *> SemaHLSL::inferGlobalBinding(
Expr *E) {
5506 if (
auto *Ternary = dyn_cast<ConditionalOperator>(E)) {
5507 auto TrueInfo = inferGlobalBinding(Ternary->getTrueExpr());
5508 auto FalseInfo = inferGlobalBinding(Ternary->getFalseExpr());
5509 if (!TrueInfo || !FalseInfo)
5510 return std::nullopt;
5511 if (*TrueInfo != *FalseInfo)
5512 return std::nullopt;
5516 if (
auto *ASE = dyn_cast<ArraySubscriptExpr>(E))
5525 if (
const auto *AttrResType =
5526 HLSLAttributedResourceType::findHandleTypeOnResource(Ty)) {
5528 return Bindings.getDeclBindingInfo(VD, RC);
5535void SemaHLSL::trackLocalResource(
VarDecl *VD,
Expr *E) {
5536 std::optional<const DeclBindingInfo *> ExprBinding = inferGlobalBinding(E);
5539 diag::warn_hlsl_assigning_local_resource_is_not_unique)
5544 if (*ExprBinding ==
nullptr)
5547 auto PrevBinding = Assigns.find(VD);
5548 if (PrevBinding == Assigns.end()) {
5550 Assigns.insert({VD, *ExprBinding});
5555 if (*ExprBinding != PrevBinding->second) {
5557 diag::warn_hlsl_assigning_local_resource_is_not_unique)
5559 SemaRef.Diag(VD->getLocation(), diag::note_var_declared_here) << VD;
5570 "expected LHS to be a resource record or array of resource records");
5571 if (Opc != BO_Assign)
5576 while (
auto *ASE = dyn_cast<ArraySubscriptExpr>(E))
5584 SemaRef.Diag(Loc, diag::err_hlsl_assign_to_global_resource) << VD;
5589 trackLocalResource(VD, RHSExpr);
5597void SemaHLSL::collectResourceBindingsOnVarDecl(
VarDecl *VD) {
5599 "expected global variable that contains HLSL resource");
5602 if (
const HLSLBufferDecl *CBufferOrTBuffer = dyn_cast<HLSLBufferDecl>(VD)) {
5603 Bindings.addDeclBindingInfo(VD, CBufferOrTBuffer->isCBuffer()
5604 ? ResourceClass::CBuffer
5605 : ResourceClass::SRV);
5618 if (
const HLSLAttributedResourceType *AttrResType =
5619 HLSLAttributedResourceType::findHandleTypeOnResource(Ty)) {
5620 Bindings.addDeclBindingInfo(VD, AttrResType->getAttrs().ResourceClass);
5625 if (
const RecordType *RT = dyn_cast<RecordType>(Ty))
5626 collectResourceBindingsOnUserRecordDecl(VD, RT);
5632void SemaHLSL::processExplicitBindingsOnDecl(
VarDecl *VD) {
5635 bool HasBinding =
false;
5636 for (Attr *A : VD->
attrs()) {
5639 if (
auto PA = VD->
getAttr<HLSLVkPushConstantAttr>())
5640 Diag(PA->getLoc(), diag::err_hlsl_attr_incompatible) << A << PA;
5643 HLSLResourceBindingAttr *RBA = dyn_cast<HLSLResourceBindingAttr>(A);
5644 if (!RBA || !RBA->hasRegisterSlot())
5649 assert(RT != RegisterType::I &&
"invalid or obsolete register type should "
5650 "never have an attribute created");
5652 if (RT == RegisterType::C) {
5653 if (Bindings.hasBindingInfoForDecl(VD))
5655 diag::warn_hlsl_user_defined_type_missing_member)
5656 <<
static_cast<int>(RT);
5664 if (DeclBindingInfo *BI = Bindings.getDeclBindingInfo(VD, RC)) {
5669 diag::warn_hlsl_user_defined_type_missing_member)
5670 <<
static_cast<int>(RT);
5678class InitListTransformer {
5682 QualType *DstIt =
nullptr;
5683 Expr **ArgIt =
nullptr;
5689 bool castInitializer(Expr *E) {
5690 assert(DstIt &&
"This should always be something!");
5691 if (DstIt == DestTypes.end()) {
5693 ArgExprs.push_back(E);
5698 DstIt = DestTypes.begin();
5701 Ctx, *DstIt,
false);
5706 ArgExprs.push_back(
Init);
5711 bool buildInitializerListImpl(Expr *E) {
5713 if (
auto *
Init = dyn_cast<InitListExpr>(E)) {
5714 for (
auto *SubInit :
Init->inits())
5715 if (!buildInitializerListImpl(SubInit))
5725 return castInitializer(E);
5739 if (
auto *VecTy = Ty->
getAs<VectorType>()) {
5744 for (uint64_t I = 0; I <
Size; ++I) {
5746 SizeTy, SourceLocation());
5752 if (!castInitializer(ElExpr.
get()))
5757 if (
auto *MTy = Ty->
getAs<ConstantMatrixType>()) {
5758 unsigned Rows = MTy->getNumRows();
5759 unsigned Cols = MTy->getNumColumns();
5760 QualType ElemTy = MTy->getElementType();
5762 for (
unsigned R = 0;
R < Rows; ++
R) {
5763 for (
unsigned C = 0;
C < Cols; ++
C) {
5776 if (!castInitializer(ElExpr.
get()))
5784 if (
auto *ArrTy = dyn_cast<ConstantArrayType>(Ty.
getTypePtr())) {
5788 for (uint64_t I = 0; I <
Size; ++I) {
5790 SizeTy, SourceLocation());
5795 if (!buildInitializerListImpl(ElExpr.
get()))
5802 llvm::SmallVector<CXXRecordDecl *> RecordDecls;
5803 RecordDecls.push_back(RD);
5804 while (RecordDecls.back()->getNumBases()) {
5805 CXXRecordDecl *D = RecordDecls.back();
5807 "HLSL doesn't support multiple inheritance");
5808 RecordDecls.push_back(
5811 while (!RecordDecls.empty()) {
5812 CXXRecordDecl *RD = RecordDecls.pop_back_val();
5813 for (
auto *FD : RD->
fields()) {
5814 if (FD->isUnnamedBitField())
5822 if (!buildInitializerListImpl(Res.
get()))
5830 Expr *generateInitListsImpl(QualType Ty) {
5832 assert(ArgIt != ArgExprs.end() &&
"Something is off in iteration!");
5837 llvm::SmallVector<Expr *>
Inits;
5842 if (
auto *ATy = Ty->
getAs<VectorType>()) {
5843 ElTy = ATy->getElementType();
5844 Size = ATy->getNumElements();
5845 }
else if (
auto *CMTy = Ty->
getAs<ConstantMatrixType>()) {
5846 ElTy = CMTy->getElementType();
5847 Size = CMTy->getNumElementsFlattened();
5850 ElTy = VTy->getElementType();
5851 Size = VTy->getZExtSize();
5853 for (uint64_t I = 0; I <
Size; ++I)
5854 Inits.push_back(generateInitListsImpl(ElTy));
5857 llvm::SmallVector<CXXRecordDecl *> RecordDecls;
5858 RecordDecls.push_back(RD);
5859 while (RecordDecls.back()->getNumBases()) {
5860 CXXRecordDecl *D = RecordDecls.back();
5862 "HLSL doesn't support multiple inheritance");
5863 RecordDecls.push_back(
5866 while (!RecordDecls.empty()) {
5867 CXXRecordDecl *RD = RecordDecls.pop_back_val();
5868 for (
auto *FD : RD->
fields())
5869 if (!FD->isUnnamedBitField())
5873 auto *NewInit =
new (Ctx) InitListExpr(Ctx,
Inits.front()->getBeginLoc(),
5875 NewInit->setType(Ty);
5880 llvm::SmallVector<QualType, 16> DestTypes;
5881 llvm::SmallVector<Expr *, 16> ArgExprs;
5882 InitListTransformer(Sema &SemaRef,
const InitializedEntity &Entity)
5883 : S(SemaRef), Ctx(SemaRef.getASTContext()),
5884 Wrap(Entity.
getType()->isIncompleteArrayType()) {
5885 InitTy = Entity.
getType().getNonReferenceType();
5895 DstIt = DestTypes.begin();
5898 bool buildInitializerList(Expr *E) {
return buildInitializerListImpl(E); }
5900 Expr *generateInitLists() {
5901 assert(!ArgExprs.empty() &&
5902 "Call buildInitializerList to generate argument expressions.");
5903 ArgIt = ArgExprs.begin();
5905 return generateInitListsImpl(InitTy);
5906 llvm::SmallVector<Expr *>
Inits;
5907 while (ArgIt != ArgExprs.end())
5908 Inits.push_back(generateInitListsImpl(InitTy));
5910 auto *NewInit =
new (Ctx) InitListExpr(Ctx,
Inits.front()->getBeginLoc(),
5912 llvm::APInt ArySize(64,
Inits.size());
5914 ArraySizeModifier::Normal, 0));
5926 if (
const ArrayType *AT = dyn_cast<ArrayType>(Ty)) {
5933 if (
const auto *RT = Ty->
getAs<RecordType>()) {
5937 if (
const auto *CXXRD = dyn_cast<CXXRecordDecl>(RD)) {
5957 if (
Init->getType()->isScalarType())
5960 InitListTransformer ILT(
SemaRef, Entity);
5962 for (
unsigned I = 0; I <
Init->getNumInits(); ++I) {
5970 Init->setInit(I, E);
5972 if (!ILT.buildInitializerList(E))
5975 size_t ExpectedSize = ILT.DestTypes.size();
5976 size_t ActualSize = ILT.ArgExprs.size();
5977 if (ExpectedSize == 0 && ActualSize == 0)
5984 InitTy =
SemaRef.getASTContext().removeAddrSpaceQualType(InitTy);
5986 SemaRef.Diag(
Init->getBeginLoc(), diag::err_hlsl_incorrect_num_initializers)
5987 << (int)(ExpectedSize < ActualSize) << InitTy
5988 << ExpectedSize << ActualSize;
5998 assert(ExpectedSize > 0 &&
5999 "The expected size of an incomplete array type must be at least 1.");
6001 ((ActualSize + ExpectedSize - 1) / ExpectedSize) * ExpectedSize;
6009 InitTy =
SemaRef.getASTContext().removeAddrSpaceQualType(InitTy);
6010 if (ExpectedSize != ActualSize) {
6011 int TooManyOrFew = ActualSize > ExpectedSize ? 1 : 0;
6012 SemaRef.Diag(
Init->getBeginLoc(), diag::err_hlsl_incorrect_num_initializers)
6013 << TooManyOrFew << InitTy << ExpectedSize << ActualSize;
6020 Init->resizeInits(Ctx, NewInit->getNumInits());
6021 for (
unsigned I = 0; I < NewInit->getNumInits(); ++I)
6022 Init->updateInit(Ctx, I, NewInit->getInit(I));
6030 S.
Diag(OpLoc, diag::err_builtin_matrix_invalid_member)
6040 StringRef AccessorName = CompName->
getName();
6041 assert(!AccessorName.empty() &&
"Matrix Accessor must have a name");
6043 unsigned Rows = MT->getNumRows();
6044 unsigned Cols = MT->getNumColumns();
6045 bool IsZeroBasedAccessor =
false;
6046 unsigned ChunkLen = 0;
6047 if (AccessorName.size() < 2)
6049 "length 4 for zero based: \'_mRC\' or "
6050 "length 3 for one-based: \'_RC\' accessor",
6053 if (AccessorName[0] ==
'_') {
6054 if (AccessorName[1] ==
'm') {
6055 IsZeroBasedAccessor =
true;
6062 S, AccessorName,
"zero based: \'_mRC\' or one-based: \'_RC\' accessor",
6065 if (AccessorName.size() % ChunkLen != 0) {
6066 const llvm::StringRef
Expected = IsZeroBasedAccessor
6067 ?
"zero based: '_mRC' accessor"
6068 :
"one-based: '_RC' accessor";
6073 auto isDigit = [](
char c) {
return c >=
'0' &&
c <=
'9'; };
6074 auto isZeroBasedIndex = [](
unsigned i) {
return i <= 3; };
6075 auto isOneBasedIndex = [](
unsigned i) {
return i >= 1 && i <= 4; };
6077 bool HasRepeated =
false;
6079 unsigned NumComponents = 0;
6080 const char *Begin = AccessorName.data();
6082 for (
unsigned I = 0, E = AccessorName.size(); I < E; I += ChunkLen) {
6083 const char *Chunk = Begin + I;
6084 char RowChar = 0, ColChar = 0;
6085 if (IsZeroBasedAccessor) {
6087 if (Chunk[0] !=
'_' || Chunk[1] !=
'm') {
6088 char Bad = (Chunk[0] !=
'_') ? Chunk[0] : Chunk[1];
6090 S, StringRef(&Bad, 1),
"\'_m\' prefix",
6097 if (Chunk[0] !=
'_')
6099 S, StringRef(&Chunk[0], 1),
"\'_\' prefix",
6106 bool IsDigitsError =
false;
6108 unsigned BadPos = IsZeroBasedAccessor ? 2 : 1;
6112 IsDigitsError =
true;
6116 unsigned BadPos = IsZeroBasedAccessor ? 3 : 2;
6120 IsDigitsError =
true;
6125 unsigned Row = RowChar -
'0';
6126 unsigned Col = ColChar -
'0';
6128 bool HasIndexingError =
false;
6129 if (IsZeroBasedAccessor) {
6131 if (!isZeroBasedIndex(Row)) {
6132 S.
Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
6134 HasIndexingError =
true;
6136 if (!isZeroBasedIndex(Col)) {
6137 S.
Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
6139 HasIndexingError =
true;
6143 if (!isOneBasedIndex(Row)) {
6144 S.
Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
6146 HasIndexingError =
true;
6148 if (!isOneBasedIndex(Col)) {
6149 S.
Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
6151 HasIndexingError =
true;
6158 if (HasIndexingError)
6164 bool HasBoundsError =
false;
6166 Diag(OpLoc, diag::err_hlsl_matrix_index_out_of_bounds)
6168 HasBoundsError =
true;
6171 Diag(OpLoc, diag::err_hlsl_matrix_index_out_of_bounds)
6173 HasBoundsError =
true;
6178 unsigned FlatIndex = Row * Cols + Col;
6179 if (Seen[FlatIndex])
6181 Seen[FlatIndex] =
true;
6184 if (NumComponents == 0 || NumComponents > 4) {
6185 S.
Diag(OpLoc, diag::err_hlsl_matrix_swizzle_invalid_length)
6190 QualType ElemTy = MT->getElementType();
6191 if (NumComponents == 1)
6197 for (Sema::ExtVectorDeclsType::iterator
6201 if ((*I)->getUnderlyingType() == VT)
6212 trackLocalResource(VDecl,
Init);
6214 const HLSLVkConstantIdAttr *ConstIdAttr =
6215 VDecl->
getAttr<HLSLVkConstantIdAttr>();
6222 if (!
Init->isCXX11ConstantExpr(Context, &InitValue)) {
6232 int ConstantID = ConstIdAttr->getId();
6233 llvm::APInt IDVal(Context.getIntWidth(Context.IntTy), ConstantID);
6235 ConstIdAttr->getLocation());
6239 if (
C->getType()->getCanonicalTypeUnqualified() !=
6243 Context.getTrivialTypeSourceInfo(
6244 Init->getType(),
Init->getExprLoc()),
6263 if (!Params || Params->
size() != 1)
6276 if (
auto *TTP = dyn_cast<TemplateTypeParmDecl>(P)) {
6277 if (TTP->hasDefaultArgument()) {
6278 TemplateArgs.
addArgument(TTP->getDefaultArgument());
6281 }
else if (
auto *NTTP = dyn_cast<NonTypeTemplateParmDecl>(P)) {
6282 if (NTTP->hasDefaultArgument()) {
6283 TemplateArgs.
addArgument(NTTP->getDefaultArgument());
6286 }
else if (
auto *TTPD = dyn_cast<TemplateTemplateParmDecl>(P)) {
6287 if (TTPD->hasDefaultArgument()) {
6288 TemplateArgs.
addArgument(TTPD->getDefaultArgument());
6295 return SemaRef.CheckTemplateIdType(
6297 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 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_range fields() const
bool hasBindingInfoForDecl(const VarDecl *VD) const
DeclBindingInfo * getDeclBindingInfo(const VarDecl *VD, ResourceClass ResClass)
DeclBindingInfo * addDeclBindingInfo(const VarDecl *VD, ResourceClass ResClass)
Scope - A scope is a transient data structure that is used while parsing the program.
ASTContext & getASTContext() const
SemaDiagnosticBuilder Diag(SourceLocation Loc, unsigned DiagID)
Emit a diagnostic.
ExprResult ActOnOutParamExpr(ParmVarDecl *Param, Expr *Arg)
HLSLRootSignatureDecl * lookupRootSignatureOverrideDecl(DeclContext *DC) const
bool CanPerformElementwiseCast(Expr *Src, QualType DestType)
void handleWaveSizeAttr(Decl *D, const ParsedAttr &AL)
void handleVkLocationAttr(Decl *D, const ParsedAttr &AL)
HLSLAttributedResourceLocInfo TakeLocForHLSLAttribute(const HLSLAttributedResourceType *RT)
void handleSemanticAttr(Decl *D, const ParsedAttr &AL)
bool CanPerformScalarCast(QualType SrcTy, QualType DestTy)
QualType ProcessResourceTypeAttributes(QualType Wrapped)
void handleShaderAttr(Decl *D, const ParsedAttr &AL)
uint32_t getNextImplicitBindingOrderID()
void CheckEntryPoint(FunctionDecl *FD)
void handleVkExtBuiltinOutputAttr(Decl *D, const ParsedAttr &AL)
void emitLogicalOperatorFixIt(Expr *LHS, Expr *RHS, BinaryOperatorKind Opc)
T * createSemanticAttr(const AttributeCommonInfo &ACI, std::optional< unsigned > Location)
bool initGlobalResourceDecl(VarDecl *VD)
void ActOnEndOfTranslationUnit(TranslationUnitDecl *TU)
bool initGlobalResourceArrayDecl(VarDecl *VD)
HLSLVkConstantIdAttr * mergeVkConstantIdAttr(Decl *D, const AttributeCommonInfo &AL, int Id)
HLSLNumThreadsAttr * mergeNumThreadsAttr(Decl *D, const AttributeCommonInfo &AL, int X, int Y, int Z)
void deduceAddressSpace(VarDecl *Decl)
std::pair< IdentifierInfo *, bool > ActOnStartRootSignatureDecl(StringRef Signature)
Computes the unique Root Signature identifier from the given signature, then lookup if there is a pre...
void handlePackOffsetAttr(Decl *D, const ParsedAttr &AL)
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 TemporaryMaterializationConversion(Expr *E)
If E is a prvalue denoting an unmaterialized temporary, materialize it as an xvalue.
ExprResult BuildFieldReferenceExpr(Expr *BaseExpr, bool IsArrow, SourceLocation OpLoc, const CXXScopeSpec &SS, FieldDecl *Field, DeclAccessPair FoundDecl, const DeclarationNameInfo &MemberNameInfo)
bool checkArgCountRange(CallExpr *Call, unsigned MinArgCount, unsigned MaxArgCount)
Checks that a call expression's argument count is in the desired range.
ExternalSemaSource * getExternalSource() const
bool checkArgCount(CallExpr *Call, unsigned DesiredArgCount)
Checks that a call expression's argument count is the desired number.
ExprResult CreateBuiltinArraySubscriptExpr(Expr *Base, SourceLocation LLoc, Expr *Idx, SourceLocation RLoc)
bool LookupQualifiedName(LookupResult &R, DeclContext *LookupCtx, bool InUnqualifiedLookup=false)
Perform qualified name lookup into a given context.
ExprResult PerformCopyInitialization(const InitializedEntity &Entity, SourceLocation EqualLoc, ExprResult Init, bool TopLevelOfInitList=false, bool AllowExplicit=false)
ExprResult CreateBuiltinMatrixSubscriptExpr(Expr *Base, Expr *RowIdx, Expr *ColumnIdx, SourceLocation RBLoc)
Encodes a location in the source.
SourceLocation getLocWithOffset(IntTy Offset) const
Return a source location with the specified offset from this SourceLocation.
A trivial tuple used to represent a source range.
SourceLocation getEnd() const
SourceLocation getEndLoc() const LLVM_READONLY
void printPretty(raw_ostream &OS, PrinterHelper *Helper, const PrintingPolicy &Policy, unsigned Indentation=0, StringRef NewlineSymbol="\n", const ASTContext *Context=nullptr) const
SourceRange getSourceRange() const LLVM_READONLY
SourceLocation tokens are not useful in isolation - they are low level value objects created/interpre...
SourceLocation getBeginLoc() const LLVM_READONLY
StringLiteral - This represents a string literal expression, e.g.
static StringLiteral * Create(const ASTContext &Ctx, StringRef Str, StringLiteralKind Kind, bool Pascal, QualType Ty, ArrayRef< SourceLocation > Locs)
This is the "fully general" constructor that allows representation of strings formed from one or more...
void startDefinition()
Starts the definition of this tag declaration.
Exposes information about the current target.
TargetOptions & getTargetOpts() const
Retrieve the target options.
const llvm::Triple & getTriple() const
Returns the target triple of the primary target.
StringRef getPlatformName() const
Retrieve the name of the platform as it is used in the availability attribute.
VersionTuple getPlatformMinVersion() const
Retrieve the minimum desired version of the platform, to which the program should be compiled.
std::string HLSLEntry
The entry point name for HLSL shader being compiled as specified by -E.
A convenient class for passing around template argument information.
void addArgument(const TemplateArgumentLoc &Loc)
The base class of all kinds of template declarations (e.g., class, function, etc.).
Stores a list of template parameters for a TemplateDecl and its derived classes.
The top declaration context.
SourceLocation getBeginLoc() const
Get the begin source location.
A container of type source information.
TypeLoc getTypeLoc() const
Return the TypeLoc wrapper for the type source info.
The base class of the type hierarchy.
bool isBooleanType() const
bool isIncompleteArrayType() const
CXXRecordDecl * getAsCXXRecordDecl() const
Retrieves the CXXRecordDecl that this type refers to, either because the type is a RecordType or beca...
bool isConstantArrayType() const
bool hasIntegerRepresentation() const
Determine whether this type has an integer representation of some sort, e.g., it is an integer type o...
CXXRecordDecl * castAsCXXRecordDecl() const
bool isArithmeticType() const
bool isConstantMatrixType() const
bool isHLSLBuiltinIntangibleType() const
bool isPointerType() const
CanQualType getCanonicalTypeUnqualified() const
bool isIntegerType() const
isIntegerType() does not include complex integers (a GCC extension).
const T * castAs() const
Member-template castAs<specific type>.
bool isReferenceType() const
bool isHLSLIntangibleType() const
bool isEnumeralType() const
bool isScalarType() const
bool isIntegralType(const ASTContext &Ctx) const
Determine whether this type is an integral type.
const Type * getArrayElementTypeNoTypeQual() const
If this is an array type, return the element type of the array, potentially with type qualifiers miss...
QualType getPointeeType() const
If this is a pointer, ObjC object pointer, or block pointer, this returns the respective pointee.
bool hasUnsignedIntegerRepresentation() const
Determine whether this type has an unsigned integer representation of some sort, e....
bool isAggregateType() const
Determines whether the type is a C++ aggregate type or C aggregate or union type.
ScalarTypeKind getScalarTypeKind() const
Given that this is a scalar type, classify it.
bool hasSignedIntegerRepresentation() const
Determine whether this type has an signed integer representation of some sort, e.g....
bool isMatrixType() const
bool isHLSLResourceRecord() const
bool hasFloatingRepresentation() const
Determine whether this type has a floating-point representation of some sort, e.g....
bool isVectorType() const
bool isRealFloatingType() const
Floating point categories.
bool isHLSLAttributedResourceType() const
bool isFloatingType() const
const T * getAs() const
Member-template getAs<specific type>'.
const Type * getUnqualifiedDesugaredType() const
Return the specified type with any "sugar" removed from the type, removing any typedefs,...
bool isRecordType() const
bool isHLSLResourceRecordArray() const
void setType(QualType newType)
Represents a variable declaration or definition.
static VarDecl * Create(ASTContext &C, DeclContext *DC, SourceLocation StartLoc, SourceLocation IdLoc, const IdentifierInfo *Id, QualType T, TypeSourceInfo *TInfo, StorageClass S)
void setInitStyle(InitializationStyle Style)
@ CallInit
Call-style initialization (C++98)
void setStorageClass(StorageClass SC)
bool hasGlobalStorage() const
Returns true for all variables that do not have local storage.
StorageClass getStorageClass() const
Returns the storage class as written in the source.
Represents a GCC generic vector type.
unsigned getNumElements() const
QualType getElementType() const
void pushName(llvm::StringRef N)
void pushArrayIndex(uint64_t Index)
void pushBaseName(llvm::StringRef N)
IdentifierInfo * getNameAsIdentifier(ASTContext &AST) const
Defines the clang::TargetInfo interface.
uint32_t getResourceDimensions(llvm::dxil::ResourceDimension Dim)
bool hasCounterHandle(const CXXRecordDecl *RD)
The JSON file list parser is used to communicate input to InstallAPI.
bool isa(CodeGen::Address addr)
if(T->getSizeExpr()) TRY_TO(TraverseStmt(const_cast< Expr * >(T -> getSizeExpr())))
static bool CheckFloatOrHalfRepresentation(Sema *S, SourceLocation Loc, int ArgOrdinal, clang::QualType PassedType)
@ ICIS_NoInit
No in-class initializer.
@ TemplateName
The identifier is a template name. FIXME: Add an annotation for that.
@ OK_Ordinary
An ordinary object is located at an address in memory.
static bool CheckAllArgTypesAreCorrect(Sema *S, CallExpr *TheCall, llvm::ArrayRef< llvm::function_ref< bool(Sema *, SourceLocation, int, QualType)> > Checks)
@ AANT_ArgumentIdentifier
@ Result
The result type of a method or function.
@ Ordinary
This parameter uses ordinary ABI rules for its type.
llvm::Expected< QualType > ExpectedType
@ Template
We are parsing a template declaration.
LLVM_READONLY bool isDigit(unsigned char c)
Return true if this character is an ASCII digit: [0-9].
static bool CheckAllArgsHaveSameType(Sema *S, CallExpr *TheCall)
@ Type
The name was classified as a type.
LangAS
Defines the address space values used by the address space qualifier of QualType.
bool CreateHLSLAttributedResourceType(Sema &S, QualType Wrapped, ArrayRef< const Attr * > AttrList, QualType &ResType, HLSLAttributedResourceLocInfo *LocInfo=nullptr)
CastKind
CastKind - The kind of operation required for a conversion.
ExprValueKind
The categorization of expression values, currently following the C++11 scheme.
@ VK_PRValue
A pr-value expression (in the C++11 taxonomy) produces a temporary value.
@ VK_LValue
An l-value expression is a reference to an object with independent storage.
DynamicRecursiveASTVisitorBase< false > DynamicRecursiveASTVisitor
U cast(CodeGen::Address addr)
@ None
No keyword precedes the qualified type name.
ActionResult< Expr * > ExprResult
Visibility
Describes the different kinds of visibility that a declaration may have.
hash_code hash_value(const clang::dependencies::ModuleID &ID)
__DEVICE__ bool isnan(float __x)
__DEVICE__ _Tp abs(const std::complex< _Tp > &__c)
TypeSourceInfo * ContainedTyInfo
Describes how types, statements, expressions, and declarations should be printed.
unsigned getImplicitOrderID() const
void setCounterImplicitOrderID(unsigned Value) const
bool hasCounterImplicitOrderID() const
unsigned getSpace() const
bool hasImplicitOrderID() const
void setImplicitOrderID(unsigned Value) const
const SourceLocation & getLocation() const
const llvm::hlsl::rootsig::RootElement & getElement() const