37#include "llvm/ADT/ArrayRef.h"
38#include "llvm/ADT/STLExtras.h"
39#include "llvm/ADT/SmallVector.h"
40#include "llvm/ADT/StringExtras.h"
41#include "llvm/ADT/StringRef.h"
42#include "llvm/ADT/Twine.h"
43#include "llvm/Frontend/HLSL/HLSLBinding.h"
44#include "llvm/Frontend/HLSL/RootSignatureValidations.h"
45#include "llvm/Support/Casting.h"
46#include "llvm/Support/DXILABI.h"
47#include "llvm/Support/ErrorHandling.h"
48#include "llvm/Support/FormatVariadic.h"
49#include "llvm/TargetParser/Triple.h"
64 case ResourceClass::SRV:
65 return RegisterType::SRV;
66 case ResourceClass::UAV:
67 return RegisterType::UAV;
68 case ResourceClass::CBuffer:
69 return RegisterType::CBuffer;
70 case ResourceClass::Sampler:
71 return RegisterType::Sampler;
73 llvm_unreachable(
"unexpected ResourceClass value");
82 case ResourceClass::SRV:
83 case ResourceClass::UAV:
85 case ResourceClass::CBuffer:
87 case ResourceClass::Sampler:
90 llvm_unreachable(
"unexpected ResourceClass value");
96 assert(RT !=
nullptr);
100 *RT = RegisterType::SRV;
104 *RT = RegisterType::UAV;
108 *RT = RegisterType::CBuffer;
112 *RT = RegisterType::Sampler;
116 *RT = RegisterType::C;
120 *RT = RegisterType::I;
129 case RegisterType::SRV:
131 case RegisterType::UAV:
133 case RegisterType::CBuffer:
135 case RegisterType::Sampler:
137 case RegisterType::C:
139 case RegisterType::I:
142 llvm_unreachable(
"unexpected RegisterType value");
147 case RegisterType::SRV:
148 return ResourceClass::SRV;
149 case RegisterType::UAV:
150 return ResourceClass::UAV;
151 case RegisterType::CBuffer:
152 return ResourceClass::CBuffer;
153 case RegisterType::Sampler:
154 return ResourceClass::Sampler;
155 case RegisterType::C:
156 case RegisterType::I:
160 llvm_unreachable(
"unexpected RegisterType value");
164 const auto *BT = dyn_cast<BuiltinType>(
Type);
168 return Builtin::BI__builtin_get_spirv_spec_constant_int;
171 switch (BT->getKind()) {
172 case BuiltinType::Bool:
173 return Builtin::BI__builtin_get_spirv_spec_constant_bool;
174 case BuiltinType::Short:
175 return Builtin::BI__builtin_get_spirv_spec_constant_short;
176 case BuiltinType::Int:
177 return Builtin::BI__builtin_get_spirv_spec_constant_int;
178 case BuiltinType::LongLong:
179 return Builtin::BI__builtin_get_spirv_spec_constant_longlong;
180 case BuiltinType::UShort:
181 return Builtin::BI__builtin_get_spirv_spec_constant_ushort;
182 case BuiltinType::UInt:
183 return Builtin::BI__builtin_get_spirv_spec_constant_uint;
184 case BuiltinType::ULongLong:
185 return Builtin::BI__builtin_get_spirv_spec_constant_ulonglong;
186 case BuiltinType::Half:
187 return Builtin::BI__builtin_get_spirv_spec_constant_half;
188 case BuiltinType::Float:
189 return Builtin::BI__builtin_get_spirv_spec_constant_float;
190 case BuiltinType::Double:
191 return Builtin::BI__builtin_get_spirv_spec_constant_double;
200 llvm::raw_svector_ostream OS(Buffer);
207 ResourceClass ResClass) {
209 "DeclBindingInfo already added");
215 DeclToBindingListIndex.try_emplace(VD, BindingsList.size());
216 return &BindingsList.emplace_back(VD, ResClass);
220 ResourceClass ResClass) {
221 auto Entry = DeclToBindingListIndex.find(VD);
222 if (Entry != DeclToBindingListIndex.end()) {
223 for (
unsigned Index = Entry->getSecond();
224 Index < BindingsList.size() && BindingsList[Index].Decl == VD;
226 if (BindingsList[Index].ResClass == ResClass)
227 return &BindingsList[Index];
234 return DeclToBindingListIndex.contains(VD);
246 getASTContext(), LexicalParent, CBuffer, KwLoc, Ident, IdentLoc, LBrace);
249 auto RC = CBuffer ? llvm::hlsl::ResourceClass::CBuffer
250 : llvm::hlsl::ResourceClass::SRV;
262 if (T->isArrayType() || T->isStructureType() || T->isConstantMatrixType())
269 assert(Context.getTypeSize(T) <= 64 &&
270 "Scalar bit widths larger than 64 not supported");
273 return Context.getTypeSize(T) / 8;
280 constexpr unsigned CBufferAlign = 16;
281 if (
const auto *RD = T->getAsRecordDecl()) {
283 for (
const FieldDecl *Field : RD->fields()) {
290 unsigned AlignSize = llvm::alignTo(Size, FieldAlign);
291 if ((AlignSize % CBufferAlign) + FieldSize > CBufferAlign) {
292 FieldAlign = CBufferAlign;
295 Size = llvm::alignTo(Size, FieldAlign);
302 unsigned ElementCount = AT->getSize().getZExtValue();
303 if (ElementCount == 0)
306 unsigned ElementSize =
308 unsigned AlignedElementSize = llvm::alignTo(ElementSize, CBufferAlign);
309 return AlignedElementSize * (ElementCount - 1) + ElementSize;
313 unsigned ElementCount = VT->getNumElements();
314 unsigned ElementSize =
316 return ElementSize * ElementCount;
319 return Context.getTypeSize(T) / 8;
330 bool HasPackOffset =
false;
331 bool HasNonPackOffset =
false;
333 VarDecl *Var = dyn_cast<VarDecl>(Field);
336 if (Field->hasAttr<HLSLPackOffsetAttr>()) {
337 PackOffsetVec.emplace_back(Var, Field->
getAttr<HLSLPackOffsetAttr>());
338 HasPackOffset =
true;
340 HasNonPackOffset =
true;
347 if (HasNonPackOffset)
354 std::sort(PackOffsetVec.begin(), PackOffsetVec.end(),
355 [](
const std::pair<VarDecl *, HLSLPackOffsetAttr *> &LHS,
356 const std::pair<VarDecl *, HLSLPackOffsetAttr *> &RHS) {
357 return LHS.second->getOffsetInBytes() <
358 RHS.second->getOffsetInBytes();
360 for (
unsigned i = 0; i < PackOffsetVec.size() - 1; i++) {
361 VarDecl *Var = PackOffsetVec[i].first;
362 HLSLPackOffsetAttr *
Attr = PackOffsetVec[i].second;
364 unsigned Begin =
Attr->getOffsetInBytes();
365 unsigned End = Begin + Size;
366 unsigned NextBegin = PackOffsetVec[i + 1].second->getOffsetInBytes();
367 if (End > NextBegin) {
368 VarDecl *NextVar = PackOffsetVec[i + 1].first;
380 CAT = dyn_cast<ConstantArrayType>(
382 return CAT !=
nullptr;
393static const HLSLAttributedResourceType *
396 "expected array of resource records");
398 while (
const ArrayType *AT = dyn_cast<ArrayType>(Ty))
400 return HLSLAttributedResourceType::findHandleTypeOnResource(Ty);
403static const HLSLAttributedResourceType *
417 return RD->isEmpty();
446 Base.getType()->castAsCXXRecordDecl()))
457 assert(RD ==
nullptr &&
458 "there should be at most 1 record by a given name in a scope");
475 Name.append(NameBaseII->
getName());
482 size_t NameLength = Name.size();
491 Name.append(llvm::Twine(suffix).str());
492 II = &AST.
Idents.
get(Name, tok::TokenKind::identifier);
499 Name.truncate(NameLength);
514 if (
const auto *CAT = dyn_cast<ConstantArrayType>(Ty)) {
516 S, CAT->getElementType()->getUnqualifiedDesugaredType());
521 CAT->getSizeModifier(),
522 CAT->getIndexTypeCVRQualifiers())
560 "struct is already HLSL buffer compatible");
574 LS->
addAttr(PackedAttr::CreateImplicit(AST));
578 if (
unsigned NumBases = StructDecl->
getNumBases()) {
579 assert(NumBases == 1 &&
"HLSL supports only one base type");
629 LS->
addAttr(PackedAttr::CreateImplicit(AST));
634 VarDecl *VD = dyn_cast<VarDecl>(D);
650 "host layout field for $Globals decl failed to be created");
667 uint32_t ImplicitBindingOrderID) {
669 HLSLResourceBindingAttr::CreateImplicit(S.
getASTContext(),
"",
"0", {});
670 Attr->setBinding(RT, std::nullopt, 0);
671 Attr->setImplicitBindingOrderID(ImplicitBindingOrderID);
678 BufDecl->setRBraceLoc(RBrace);
695 BufDecl->isCBuffer() ? RegisterType::CBuffer
705 int X,
int Y,
int Z) {
706 if (HLSLNumThreadsAttr *NT = D->
getAttr<HLSLNumThreadsAttr>()) {
707 if (NT->getX() !=
X || NT->getY() != Y || NT->getZ() != Z) {
708 Diag(NT->getLocation(), diag::err_hlsl_attribute_param_mismatch) << AL;
709 Diag(AL.
getLoc(), diag::note_conflicting_attribute);
719 int Min,
int Max,
int Preferred,
720 int SpelledArgsCount) {
721 if (HLSLWaveSizeAttr *WS = D->
getAttr<HLSLWaveSizeAttr>()) {
722 if (WS->getMin() !=
Min || WS->getMax() !=
Max ||
723 WS->getPreferred() != Preferred ||
724 WS->getSpelledArgsCount() != SpelledArgsCount) {
725 Diag(WS->getLocation(), diag::err_hlsl_attribute_param_mismatch) << AL;
726 Diag(AL.
getLoc(), diag::note_conflicting_attribute);
732 Result->setSpelledArgsCount(SpelledArgsCount);
736HLSLVkConstantIdAttr *
742 Diag(AL.
getLoc(), diag::warn_attribute_ignored) << AL;
750 Diag(VD->getLocation(), diag::err_specialization_const);
754 if (!VD->getType().isConstQualified()) {
755 Diag(VD->getLocation(), diag::err_specialization_const);
759 if (HLSLVkConstantIdAttr *CI = D->
getAttr<HLSLVkConstantIdAttr>()) {
760 if (CI->getId() != Id) {
761 Diag(CI->getLocation(), diag::err_hlsl_attribute_param_mismatch) << AL;
762 Diag(AL.
getLoc(), diag::note_conflicting_attribute);
767 HLSLVkConstantIdAttr *
Result =
774 llvm::Triple::EnvironmentType ShaderType) {
775 if (HLSLShaderAttr *NT = D->
getAttr<HLSLShaderAttr>()) {
776 if (NT->getType() != ShaderType) {
777 Diag(NT->getLocation(), diag::err_hlsl_attribute_param_mismatch) << AL;
778 Diag(AL.
getLoc(), diag::note_conflicting_attribute);
782 return HLSLShaderAttr::Create(
getASTContext(), ShaderType, AL);
785HLSLParamModifierAttr *
787 HLSLParamModifierAttr::Spelling Spelling) {
790 if (HLSLParamModifierAttr *PA = D->
getAttr<HLSLParamModifierAttr>()) {
791 if ((PA->isIn() && Spelling == HLSLParamModifierAttr::Keyword_out) ||
792 (PA->isOut() && Spelling == HLSLParamModifierAttr::Keyword_in)) {
793 D->
dropAttr<HLSLParamModifierAttr>();
795 return HLSLParamModifierAttr::Create(
797 HLSLParamModifierAttr::Keyword_inout);
799 Diag(AL.
getLoc(), diag::err_hlsl_duplicate_parameter_modifier) << AL;
800 Diag(PA->getLocation(), diag::note_conflicting_attribute);
826 if (HLSLShaderAttr::isValidShaderType(Env) && Env != llvm::Triple::Library) {
827 if (
const auto *Shader = FD->
getAttr<HLSLShaderAttr>()) {
830 if (Shader->getType() != Env) {
831 Diag(Shader->getLocation(), diag::err_hlsl_entry_shader_attr_mismatch)
843 case llvm::Triple::UnknownEnvironment:
844 case llvm::Triple::Library:
846 case llvm::Triple::RootSignature:
847 llvm_unreachable(
"rootsig environment has no functions");
849 llvm_unreachable(
"Unhandled environment in triple");
855 HLSLAppliedSemanticAttr *Semantic,
860 const auto *ShaderAttr = FD->
getAttr<HLSLShaderAttr>();
861 assert(ShaderAttr &&
"Entry point has no shader attribute");
862 llvm::Triple::EnvironmentType ST = ShaderAttr->getType();
863 auto SemanticName = Semantic->getSemanticName().upper();
868 if (SemanticName ==
"SV_POSITION") {
869 return (ST == llvm::Triple::Vertex && !IsInput) ||
870 (ST == llvm::Triple::Pixel && IsInput);
872 if (SemanticName ==
"SV_VERTEXID")
878bool SemaHLSL::determineActiveSemanticOnScalar(
FunctionDecl *FD,
881 SemanticInfo &ActiveSemantic,
882 SemaHLSL::SemanticContext &SC) {
883 if (ActiveSemantic.Semantic ==
nullptr) {
884 ActiveSemantic.Semantic = D->
getAttr<HLSLParsedSemanticAttr>();
885 if (ActiveSemantic.Semantic)
886 ActiveSemantic.Index = ActiveSemantic.Semantic->getSemanticIndex();
889 if (!ActiveSemantic.Semantic) {
895 HLSLAppliedSemanticAttr(
getASTContext(), *ActiveSemantic.Semantic,
896 ActiveSemantic.Semantic->getAttrName()->getName(),
897 ActiveSemantic.Index.value_or(0));
901 checkSemanticAnnotation(FD, D, A, SC);
902 OutputDecl->addAttr(A);
904 unsigned Location = ActiveSemantic.Index.value_or(0);
907 SC.CurrentIOType & IOType::In)) {
908 bool HasVkLocation =
false;
909 if (
auto *A = D->getAttr<HLSLVkLocationAttr>()) {
910 HasVkLocation = true;
911 Location = A->getLocation();
914 if (SC.UsesExplicitVkLocations.value_or(HasVkLocation) != HasVkLocation) {
915 Diag(D->getLocation(), diag::err_hlsl_semantic_partial_explicit_indexing);
918 SC.UsesExplicitVkLocations = HasVkLocation;
921 const ConstantArrayType *AT = dyn_cast<ConstantArrayType>(D->getType());
922 unsigned ElementCount = AT ? AT->
getZExtSize() : 1;
923 ActiveSemantic.Index = Location + ElementCount;
925 Twine BaseName = Twine(ActiveSemantic.Semantic->getAttrName()->getName());
926 for (
unsigned I = 0; I < ElementCount; ++I) {
927 Twine VariableName = BaseName.concat(Twine(Location + I));
929 auto [_, Inserted] = SC.ActiveSemantics.insert(VariableName.str());
931 Diag(D->getLocation(), diag::err_hlsl_semantic_index_overlap)
932 << VariableName.str();
943 SemanticInfo &ActiveSemantic,
944 SemaHLSL::SemanticContext &SC) {
945 if (ActiveSemantic.Semantic ==
nullptr) {
946 ActiveSemantic.Semantic = D->
getAttr<HLSLParsedSemanticAttr>();
947 if (ActiveSemantic.Semantic)
948 ActiveSemantic.Index = ActiveSemantic.Semantic->getSemanticIndex();
954 const RecordType *RT = dyn_cast<RecordType>(T);
956 return determineActiveSemanticOnScalar(FD, OutputDecl, D, ActiveSemantic,
959 const RecordDecl *RD = RT->getDecl();
960 for (FieldDecl *Field : RD->
fields()) {
961 SemanticInfo Info = ActiveSemantic;
962 if (!determineActiveSemantic(FD, OutputDecl, Field, Info, SC)) {
963 Diag(
Field->getLocation(), diag::note_hlsl_semantic_used_here) <<
Field;
966 if (ActiveSemantic.Semantic)
967 ActiveSemantic = Info;
974 const auto *ShaderAttr = FD->
getAttr<HLSLShaderAttr>();
975 assert(ShaderAttr &&
"Entry point has no shader attribute");
976 llvm::Triple::EnvironmentType ST = ShaderAttr->getType();
980 case llvm::Triple::Pixel:
981 case llvm::Triple::Vertex:
982 case llvm::Triple::Geometry:
983 case llvm::Triple::Hull:
984 case llvm::Triple::Domain:
985 case llvm::Triple::RayGeneration:
986 case llvm::Triple::Intersection:
987 case llvm::Triple::AnyHit:
988 case llvm::Triple::ClosestHit:
989 case llvm::Triple::Miss:
990 case llvm::Triple::Callable:
991 if (
const auto *NT = FD->
getAttr<HLSLNumThreadsAttr>()) {
992 diagnoseAttrStageMismatch(NT, ST,
993 {llvm::Triple::Compute,
994 llvm::Triple::Amplification,
995 llvm::Triple::Mesh});
998 if (
const auto *WS = FD->
getAttr<HLSLWaveSizeAttr>()) {
999 diagnoseAttrStageMismatch(WS, ST,
1000 {llvm::Triple::Compute,
1001 llvm::Triple::Amplification,
1002 llvm::Triple::Mesh});
1007 case llvm::Triple::Compute:
1008 case llvm::Triple::Amplification:
1009 case llvm::Triple::Mesh:
1010 if (!FD->
hasAttr<HLSLNumThreadsAttr>()) {
1012 << llvm::Triple::getEnvironmentTypeName(ST);
1015 if (
const auto *WS = FD->
getAttr<HLSLWaveSizeAttr>()) {
1016 if (Ver < VersionTuple(6, 6)) {
1017 Diag(WS->getLocation(), diag::err_hlsl_attribute_in_wrong_shader_model)
1020 }
else if (WS->getSpelledArgsCount() > 1 && Ver < VersionTuple(6, 8)) {
1023 diag::err_hlsl_attribute_number_arguments_insufficient_shader_model)
1024 << WS << WS->getSpelledArgsCount() <<
"6.8";
1029 case llvm::Triple::RootSignature:
1030 llvm_unreachable(
"rootsig environment has no function entry point");
1032 llvm_unreachable(
"Unhandled environment in triple");
1035 SemaHLSL::SemanticContext InputSC = {};
1036 InputSC.CurrentIOType = IOType::In;
1039 SemanticInfo ActiveSemantic;
1040 ActiveSemantic.Semantic = Param->getAttr<HLSLParsedSemanticAttr>();
1041 if (ActiveSemantic.Semantic)
1042 ActiveSemantic.Index = ActiveSemantic.Semantic->getSemanticIndex();
1045 if (!determineActiveSemantic(FD, Param, Param, ActiveSemantic, InputSC)) {
1046 Diag(Param->getLocation(), diag::note_previous_decl) << Param;
1051 SemanticInfo ActiveSemantic;
1052 SemaHLSL::SemanticContext OutputSC = {};
1053 OutputSC.CurrentIOType = IOType::Out;
1054 ActiveSemantic.Semantic = FD->
getAttr<HLSLParsedSemanticAttr>();
1055 if (ActiveSemantic.Semantic)
1056 ActiveSemantic.Index = ActiveSemantic.Semantic->getSemanticIndex();
1058 determineActiveSemantic(FD, FD, FD, ActiveSemantic, OutputSC);
1061void SemaHLSL::checkSemanticAnnotation(
1063 const HLSLAppliedSemanticAttr *SemanticAttr,
const SemanticContext &SC) {
1064 auto *ShaderAttr = EntryPoint->
getAttr<HLSLShaderAttr>();
1065 assert(ShaderAttr &&
"Entry point has no shader attribute");
1066 llvm::Triple::EnvironmentType ST = ShaderAttr->getType();
1068 auto SemanticName = SemanticAttr->getSemanticName().upper();
1069 if (SemanticName ==
"SV_DISPATCHTHREADID" ||
1070 SemanticName ==
"SV_GROUPINDEX" || SemanticName ==
"SV_GROUPTHREADID" ||
1071 SemanticName ==
"SV_GROUPID") {
1073 if (ST != llvm::Triple::Compute)
1074 diagnoseSemanticStageMismatch(SemanticAttr, ST, SC.CurrentIOType,
1075 {{llvm::Triple::Compute, IOType::In}});
1077 if (SemanticAttr->getSemanticIndex() != 0) {
1078 std::string PrettyName =
1079 "'" + SemanticAttr->getSemanticName().str() +
"'";
1080 Diag(SemanticAttr->getLoc(),
1081 diag::err_hlsl_semantic_indexing_not_supported)
1087 if (SemanticName ==
"SV_POSITION") {
1090 diagnoseSemanticStageMismatch(SemanticAttr, ST, SC.CurrentIOType,
1091 {{llvm::Triple::Vertex, IOType::InOut},
1092 {llvm::Triple::Pixel, IOType::In}});
1095 if (SemanticName ==
"SV_VERTEXID") {
1096 diagnoseSemanticStageMismatch(SemanticAttr, ST, SC.CurrentIOType,
1097 {{llvm::Triple::Vertex, IOType::In}});
1101 if (SemanticName ==
"SV_TARGET") {
1102 diagnoseSemanticStageMismatch(SemanticAttr, ST, SC.CurrentIOType,
1103 {{llvm::Triple::Pixel, IOType::Out}});
1109 if (SemanticAttr->getAttrName()->getName().starts_with_insensitive(
"SV_"))
1110 llvm_unreachable(
"Unknown SemanticAttr");
1113void SemaHLSL::diagnoseAttrStageMismatch(
1114 const Attr *A, llvm::Triple::EnvironmentType Stage,
1115 std::initializer_list<llvm::Triple::EnvironmentType> AllowedStages) {
1116 SmallVector<StringRef, 8> StageStrings;
1117 llvm::transform(AllowedStages, std::back_inserter(StageStrings),
1118 [](llvm::Triple::EnvironmentType ST) {
1120 HLSLShaderAttr::ConvertEnvironmentTypeToStr(ST));
1122 Diag(A->
getLoc(), diag::err_hlsl_attr_unsupported_in_stage)
1123 << A->
getAttrName() << llvm::Triple::getEnvironmentTypeName(Stage)
1124 << (AllowedStages.size() != 1) << join(StageStrings,
", ");
1127void SemaHLSL::diagnoseSemanticStageMismatch(
1128 const Attr *A, llvm::Triple::EnvironmentType Stage, IOType CurrentIOType,
1129 std::initializer_list<SemanticStageInfo> Allowed) {
1131 for (
auto &Case : Allowed) {
1132 if (Case.Stage != Stage)
1135 if (CurrentIOType & Case.AllowedIOTypesMask)
1138 SmallVector<std::string, 8> ValidCases;
1140 Allowed, std::back_inserter(ValidCases), [](SemanticStageInfo Case) {
1141 SmallVector<std::string, 2> ValidType;
1142 if (Case.AllowedIOTypesMask & IOType::In)
1143 ValidType.push_back(
"input");
1144 if (Case.AllowedIOTypesMask & IOType::Out)
1145 ValidType.push_back(
"output");
1147 HLSLShaderAttr::ConvertEnvironmentTypeToStr(Case.Stage)) +
1148 " " + join(ValidType,
"/");
1150 Diag(A->
getLoc(), diag::err_hlsl_semantic_unsupported_iotype_for_stage)
1151 << A->
getAttrName() << (CurrentIOType & IOType::In ?
"input" :
"output")
1152 << llvm::Triple::getEnvironmentTypeName(Case.Stage)
1153 << join(ValidCases,
", ");
1157 SmallVector<StringRef, 8> StageStrings;
1159 Allowed, std::back_inserter(StageStrings), [](SemanticStageInfo Case) {
1161 HLSLShaderAttr::ConvertEnvironmentTypeToStr(Case.Stage));
1164 Diag(A->
getLoc(), diag::err_hlsl_attr_unsupported_in_stage)
1165 << A->
getAttrName() << llvm::Triple::getEnvironmentTypeName(Stage)
1166 << (Allowed.size() != 1) << join(StageStrings,
", ");
1169template <CastKind Kind>
1172 Ty = VTy->getElementType();
1177template <CastKind Kind>
1189 if (LHSFloat && RHSFloat) {
1217 if (LHSSigned == RHSSigned) {
1218 if (IsCompAssign || IntOrder >= 0)
1226 if (IntOrder != (LHSSigned ? 1 : -1)) {
1227 if (IsCompAssign || RHSSigned)
1235 if (Ctx.getIntWidth(LElTy) != Ctx.getIntWidth(RElTy)) {
1236 if (IsCompAssign || LHSSigned)
1252 QualType ElTy = Ctx.getCorrespondingUnsignedType(LHSSigned ? LElTy : RElTy);
1253 QualType NewTy = Ctx.getExtVectorType(
1263 return CK_FloatingCast;
1265 return CK_IntegralCast;
1267 return CK_IntegralToFloating;
1269 return CK_FloatingToIntegral;
1275 bool IsCompAssign) {
1282 if (!LVecTy && IsCompAssign) {
1284 RHS =
SemaRef.ImpCastExprToType(RHS.
get(), RElTy, CK_HLSLVectorTruncation);
1286 if (Ctx.hasSameUnqualifiedType(LHSType, RHSType))
1288 RHS =
SemaRef.ImpCastExprToType(RHS.
get(), LHSType,
1293 unsigned EndSz = std::numeric_limits<unsigned>::max();
1296 LSz = EndSz = LVecTy->getNumElements();
1299 assert(EndSz != std::numeric_limits<unsigned>::max() &&
1300 "one of the above should have had a value");
1304 if (IsCompAssign && LSz != EndSz) {
1306 diag::err_hlsl_vector_compound_assignment_truncation)
1307 << LHSType << RHSType;
1313 if (!IsCompAssign && LVecTy && LVecTy->getNumElements() > EndSz)
1318 if (!IsCompAssign && !LVecTy)
1322 if (Ctx.hasSameUnqualifiedType(LHSType, RHSType))
1323 return Ctx.getCommonSugaredType(LHSType, RHSType);
1331 LElTy, RElTy, IsCompAssign);
1334 "HLSL Vectors can only contain integer or floating point types");
1336 LElTy, RElTy, IsCompAssign);
1341 assert((Opc == BO_LOr || Opc == BO_LAnd) &&
1342 "Called with non-logical operator");
1344 llvm::raw_svector_ostream OS(Buff);
1346 StringRef NewFnName = Opc == BO_LOr ?
"or" :
"and";
1347 OS << NewFnName <<
"(";
1357std::pair<IdentifierInfo *, bool>
1360 std::string IdStr =
"__hlsl_rootsig_decl_" + std::to_string(Hash);
1367 return {DeclIdent,
Found};
1378 for (
auto &RootSigElement : RootElements)
1379 Elements.push_back(RootSigElement.getElement());
1383 DeclIdent,
SemaRef.getLangOpts().HLSLRootSigVer, Elements);
1385 SignatureDecl->setImplicit();
1391 if (RootSigOverrideIdent) {
1394 if (
SemaRef.LookupQualifiedName(R, DC))
1395 return dyn_cast<HLSLRootSignatureDecl>(R.getFoundDecl());
1403struct PerVisibilityBindingChecker {
1406 std::array<llvm::hlsl::BindingInfoBuilder, 8> Builders;
1410 llvm::dxbc::ShaderVisibility Vis;
1415 PerVisibilityBindingChecker(
SemaHLSL *S) : S(S) {}
1417 void trackBinding(llvm::dxbc::ShaderVisibility
Visibility,
1418 llvm::dxil::ResourceClass RC, uint32_t Space,
1419 uint32_t LowerBound, uint32_t UpperBound,
1420 const hlsl::RootSignatureElement *Elem) {
1422 assert(BuilderIndex < Builders.size() &&
1423 "Not enough builders for visibility type");
1424 Builders[BuilderIndex].trackBinding(RC, Space, LowerBound, UpperBound,
1425 static_cast<const void *
>(Elem));
1427 static_assert(llvm::to_underlying(llvm::dxbc::ShaderVisibility::All) == 0,
1428 "'All' visibility must come first");
1429 if (
Visibility == llvm::dxbc::ShaderVisibility::All)
1430 for (
size_t I = 1, E = Builders.size(); I < E; ++I)
1431 Builders[I].trackBinding(RC, Space, LowerBound, UpperBound,
1432 static_cast<const void *
>(Elem));
1434 ElemInfoMap.push_back({Elem,
Visibility,
false});
1437 ElemInfo &
getInfo(
const hlsl::RootSignatureElement *Elem) {
1438 auto It = llvm::lower_bound(
1440 [](
const auto &LHS,
const auto &RHS) {
return LHS.Elem < RHS; });
1441 assert(It->Elem == Elem &&
"Element not in map");
1445 bool checkOverlap() {
1446 llvm::sort(ElemInfoMap, [](
const auto &LHS,
const auto &RHS) {
1447 return LHS.Elem < RHS.Elem;
1450 bool HadOverlap =
false;
1452 using llvm::hlsl::BindingInfoBuilder;
1453 auto ReportOverlap = [
this,
1454 &HadOverlap](
const BindingInfoBuilder &Builder,
1455 const llvm::hlsl::Binding &Reported) {
1459 static_cast<const hlsl::RootSignatureElement *
>(Reported.Cookie);
1460 const llvm::hlsl::Binding &
Previous = Builder.findOverlapping(Reported);
1461 const auto *PrevElem =
1462 static_cast<const hlsl::RootSignatureElement *
>(
Previous.Cookie);
1464 ElemInfo &Info =
getInfo(Elem);
1469 Info.Diagnosed =
true;
1471 ElemInfo &PrevInfo =
getInfo(PrevElem);
1472 llvm::dxbc::ShaderVisibility CommonVis =
1473 Info.Vis == llvm::dxbc::ShaderVisibility::All ? PrevInfo.Vis
1476 this->S->
Diag(Elem->
getLocation(), diag::err_hlsl_resource_range_overlap)
1477 << llvm::to_underlying(Reported.RC) << Reported.LowerBound
1478 << Reported.isUnbounded() << Reported.UpperBound
1483 this->S->
Diag(PrevElem->getLocation(),
1484 diag::note_hlsl_resource_range_here);
1487 for (BindingInfoBuilder &Builder : Builders)
1488 Builder.calculateBindingInfo(ReportOverlap);
1508 bool HadError =
false;
1509 auto ReportError = [
this, &HadError](
SourceLocation Loc, uint32_t LowerBound,
1510 uint32_t UpperBound) {
1512 this->
Diag(Loc, diag::err_hlsl_invalid_rootsig_value)
1513 << LowerBound << UpperBound;
1520 this->
Diag(Loc, diag::err_hlsl_invalid_rootsig_value)
1521 << llvm::formatv(
"{0:f}", LowerBound).sstr<6>()
1522 << llvm::formatv(
"{0:f}", UpperBound).sstr<6>();
1525 auto VerifyRegister = [ReportError](
SourceLocation Loc, uint32_t Register) {
1526 if (!llvm::hlsl::rootsig::verifyRegisterValue(Register))
1527 ReportError(Loc, 0, 0xfffffffe);
1530 auto VerifySpace = [ReportError](
SourceLocation Loc, uint32_t Space) {
1531 if (!llvm::hlsl::rootsig::verifyRegisterSpace(Space))
1532 ReportError(Loc, 0, 0xffffffef);
1535 const uint32_t Version =
1536 llvm::to_underlying(
SemaRef.getLangOpts().HLSLRootSigVer);
1537 const uint32_t VersionEnum = Version - 1;
1538 auto ReportFlagError = [
this, &HadError, VersionEnum](
SourceLocation Loc) {
1540 this->
Diag(Loc, diag::err_hlsl_invalid_rootsig_flag)
1547 const llvm::hlsl::rootsig::RootElement &Elem = RootSigElem.
getElement();
1548 if (
const auto *Descriptor =
1549 std::get_if<llvm::hlsl::rootsig::RootDescriptor>(&Elem)) {
1550 VerifyRegister(Loc, Descriptor->Reg.Number);
1551 VerifySpace(Loc, Descriptor->Space);
1553 if (!llvm::hlsl::rootsig::verifyRootDescriptorFlag(Version,
1555 ReportFlagError(Loc);
1556 }
else if (
const auto *Constants =
1557 std::get_if<llvm::hlsl::rootsig::RootConstants>(&Elem)) {
1558 VerifyRegister(Loc, Constants->Reg.Number);
1559 VerifySpace(Loc, Constants->Space);
1560 }
else if (
const auto *Sampler =
1561 std::get_if<llvm::hlsl::rootsig::StaticSampler>(&Elem)) {
1562 VerifyRegister(Loc, Sampler->Reg.Number);
1563 VerifySpace(Loc, Sampler->Space);
1566 "By construction, parseFloatParam can't produce a NaN from a "
1567 "float_literal token");
1569 if (!llvm::hlsl::rootsig::verifyMaxAnisotropy(Sampler->MaxAnisotropy))
1570 ReportError(Loc, 0, 16);
1571 if (!llvm::hlsl::rootsig::verifyMipLODBias(Sampler->MipLODBias))
1572 ReportFloatError(Loc, -16.f, 15.99f);
1573 }
else if (
const auto *Clause =
1574 std::get_if<llvm::hlsl::rootsig::DescriptorTableClause>(
1576 VerifyRegister(Loc, Clause->Reg.Number);
1577 VerifySpace(Loc, Clause->Space);
1579 if (!llvm::hlsl::rootsig::verifyNumDescriptors(Clause->NumDescriptors)) {
1583 ReportError(Loc, 1, 0xfffffffe);
1586 if (!llvm::hlsl::rootsig::verifyDescriptorRangeFlag(Version, Clause->Type,
1588 ReportFlagError(Loc);
1592 PerVisibilityBindingChecker BindingChecker(
this);
1593 SmallVector<std::pair<
const llvm::hlsl::rootsig::DescriptorTableClause *,
1598 const llvm::hlsl::rootsig::RootElement &Elem = RootSigElem.
getElement();
1599 if (
const auto *Descriptor =
1600 std::get_if<llvm::hlsl::rootsig::RootDescriptor>(&Elem)) {
1601 uint32_t LowerBound(Descriptor->Reg.Number);
1602 uint32_t UpperBound(LowerBound);
1604 BindingChecker.trackBinding(
1605 Descriptor->Visibility,
1606 static_cast<llvm::dxil::ResourceClass
>(Descriptor->Type),
1607 Descriptor->Space, LowerBound, UpperBound, &RootSigElem);
1608 }
else if (
const auto *Constants =
1609 std::get_if<llvm::hlsl::rootsig::RootConstants>(&Elem)) {
1610 uint32_t LowerBound(Constants->Reg.Number);
1611 uint32_t UpperBound(LowerBound);
1613 BindingChecker.trackBinding(
1614 Constants->Visibility, llvm::dxil::ResourceClass::CBuffer,
1615 Constants->Space, LowerBound, UpperBound, &RootSigElem);
1616 }
else if (
const auto *Sampler =
1617 std::get_if<llvm::hlsl::rootsig::StaticSampler>(&Elem)) {
1618 uint32_t LowerBound(Sampler->Reg.Number);
1619 uint32_t UpperBound(LowerBound);
1621 BindingChecker.trackBinding(
1622 Sampler->Visibility, llvm::dxil::ResourceClass::Sampler,
1623 Sampler->Space, LowerBound, UpperBound, &RootSigElem);
1624 }
else if (
const auto *Clause =
1625 std::get_if<llvm::hlsl::rootsig::DescriptorTableClause>(
1628 UnboundClauses.emplace_back(Clause, &RootSigElem);
1629 }
else if (
const auto *Table =
1630 std::get_if<llvm::hlsl::rootsig::DescriptorTable>(&Elem)) {
1631 assert(UnboundClauses.size() == Table->NumClauses &&
1632 "Number of unbound elements must match the number of clauses");
1633 bool HasAnySampler =
false;
1634 bool HasAnyNonSampler =
false;
1635 uint64_t Offset = 0;
1636 bool IsPrevUnbound =
false;
1637 for (
const auto &[Clause, ClauseElem] : UnboundClauses) {
1639 if (Clause->Type == llvm::dxil::ResourceClass::Sampler)
1640 HasAnySampler =
true;
1642 HasAnyNonSampler =
true;
1644 if (HasAnySampler && HasAnyNonSampler)
1645 Diag(Loc, diag::err_hlsl_invalid_mixed_resources);
1650 if (Clause->NumDescriptors == 0)
1654 Clause->Offset == llvm::hlsl::rootsig::DescriptorTableOffsetAppend;
1656 Offset = Clause->Offset;
1658 uint64_t RangeBound = llvm::hlsl::rootsig::computeRangeBound(
1659 Offset, Clause->NumDescriptors);
1661 if (IsPrevUnbound && IsAppending)
1662 Diag(Loc, diag::err_hlsl_appending_onto_unbound);
1663 else if (!llvm::hlsl::rootsig::verifyNoOverflowedOffset(RangeBound))
1664 Diag(Loc, diag::err_hlsl_offset_overflow) << Offset << RangeBound;
1667 Offset = RangeBound + 1;
1668 IsPrevUnbound = Clause->NumDescriptors ==
1669 llvm::hlsl::rootsig::NumDescriptorsUnbounded;
1672 uint32_t LowerBound(Clause->Reg.Number);
1673 uint32_t UpperBound = llvm::hlsl::rootsig::computeRangeBound(
1674 LowerBound, Clause->NumDescriptors);
1676 BindingChecker.trackBinding(
1678 static_cast<llvm::dxil::ResourceClass
>(Clause->Type), Clause->Space,
1679 LowerBound, UpperBound, ClauseElem);
1681 UnboundClauses.clear();
1685 return BindingChecker.checkOverlap();
1690 Diag(AL.
getLoc(), diag::err_attribute_wrong_number_arguments) << AL << 1;
1695 if (
auto *RS = D->
getAttr<RootSignatureAttr>()) {
1696 if (RS->getSignatureIdent() != Ident) {
1697 Diag(AL.
getLoc(), diag::err_disallowed_duplicate_attribute) << RS;
1701 Diag(AL.
getLoc(), diag::warn_duplicate_attribute_exact) << RS;
1707 if (
auto *SignatureDecl =
1708 dyn_cast<HLSLRootSignatureDecl>(R.getFoundDecl())) {
1715 llvm::VersionTuple SMVersion =
1720 uint32_t ZMax = 1024;
1721 uint32_t ThreadMax = 1024;
1722 if (IsDXIL && SMVersion.getMajor() <= 4) {
1725 }
else if (IsDXIL && SMVersion.getMajor() == 5) {
1735 diag::err_hlsl_numthreads_argument_oor)
1744 diag::err_hlsl_numthreads_argument_oor)
1753 diag::err_hlsl_numthreads_argument_oor)
1758 if (
X * Y * Z > ThreadMax) {
1759 Diag(AL.
getLoc(), diag::err_hlsl_numthreads_invalid) << ThreadMax;
1776 if (SpelledArgsCount == 0 || SpelledArgsCount > 3)
1784 if (SpelledArgsCount > 1 &&
1788 uint32_t Preferred = 0;
1789 if (SpelledArgsCount > 2 &&
1793 if (SpelledArgsCount > 2) {
1796 diag::err_attribute_power_of_two_in_range)
1797 << AL << llvm::dxil::MinWaveSize << llvm::dxil::MaxWaveSize
1802 if (Preferred < Min || Preferred >
Max) {
1804 diag::err_attribute_power_of_two_in_range)
1805 << AL <<
Min <<
Max << Preferred;
1808 }
else if (SpelledArgsCount > 1) {
1811 diag::err_attribute_power_of_two_in_range)
1812 << AL << llvm::dxil::MinWaveSize << llvm::dxil::MaxWaveSize <<
Max;
1816 Diag(AL.
getLoc(), diag::err_attribute_argument_invalid) << AL << 1;
1819 Diag(AL.
getLoc(), diag::warn_attr_min_eq_max) << AL;
1824 diag::err_attribute_power_of_two_in_range)
1825 << AL << llvm::dxil::MinWaveSize << llvm::dxil::MaxWaveSize <<
Min;
1830 HLSLWaveSizeAttr *NewAttr =
1867 uint32_t Binding = 0;
1891 if (!T->hasUnsignedIntegerRepresentation() ||
1892 (VT && VT->getNumElements() > 3)) {
1893 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type)
1894 << AL <<
"uint/uint2/uint3";
1903 if (!T->hasFloatingRepresentation() || (VT && VT->getNumElements() > 4)) {
1904 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type)
1905 << AL <<
"float/float1/float2/float3/float4";
1913 std::optional<unsigned> Index) {
1917 QualType ValueType = VD->getType();
1918 if (
auto *FD = dyn_cast<FunctionDecl>(D))
1921 bool IsOutput =
false;
1922 if (HLSLParamModifierAttr *MA = D->
getAttr<HLSLParamModifierAttr>()) {
1929 if (SemanticName ==
"SV_DISPATCHTHREADID") {
1932 Diag(AL.
getLoc(), diag::err_hlsl_semantic_output_not_supported) << AL;
1933 if (Index.has_value())
1934 Diag(AL.
getLoc(), diag::err_hlsl_semantic_indexing_not_supported) << AL;
1939 if (SemanticName ==
"SV_GROUPINDEX") {
1941 Diag(AL.
getLoc(), diag::err_hlsl_semantic_output_not_supported) << AL;
1942 if (Index.has_value())
1943 Diag(AL.
getLoc(), diag::err_hlsl_semantic_indexing_not_supported) << AL;
1948 if (SemanticName ==
"SV_GROUPTHREADID") {
1951 Diag(AL.
getLoc(), diag::err_hlsl_semantic_output_not_supported) << AL;
1952 if (Index.has_value())
1953 Diag(AL.
getLoc(), diag::err_hlsl_semantic_indexing_not_supported) << AL;
1958 if (SemanticName ==
"SV_GROUPID") {
1961 Diag(AL.
getLoc(), diag::err_hlsl_semantic_output_not_supported) << AL;
1962 if (Index.has_value())
1963 Diag(AL.
getLoc(), diag::err_hlsl_semantic_indexing_not_supported) << AL;
1968 if (SemanticName ==
"SV_POSITION") {
1969 const auto *VT = ValueType->getAs<
VectorType>();
1970 if (!ValueType->hasFloatingRepresentation() ||
1971 (VT && VT->getNumElements() > 4))
1972 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type)
1973 << AL <<
"float/float1/float2/float3/float4";
1978 if (SemanticName ==
"SV_VERTEXID") {
1979 uint64_t SizeInBits =
SemaRef.Context.getTypeSize(ValueType);
1980 if (!ValueType->isUnsignedIntegerType() || SizeInBits != 32)
1981 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type) << AL <<
"uint";
1986 if (SemanticName ==
"SV_TARGET") {
1987 const auto *VT = ValueType->getAs<
VectorType>();
1988 if (!ValueType->hasFloatingRepresentation() ||
1989 (VT && VT->getNumElements() > 4))
1990 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type)
1991 << AL <<
"float/float1/float2/float3/float4";
1996 Diag(AL.
getLoc(), diag::err_hlsl_unknown_semantic) << AL;
2000 uint32_t IndexValue(0), ExplicitIndex(0);
2003 assert(0 &&
"HLSLUnparsedSemantic is expected to have 2 int arguments.");
2005 assert(IndexValue > 0 ? ExplicitIndex :
true);
2006 std::optional<unsigned> Index =
2007 ExplicitIndex ? std::optional<unsigned>(IndexValue) : std::nullopt;
2017 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_ast_node)
2018 << AL <<
"shader constant in a constant buffer";
2022 uint32_t SubComponent;
2032 bool IsAggregateTy = (T->isArrayType() || T->isStructureType());
2037 if (IsAggregateTy) {
2038 Diag(AL.
getLoc(), diag::err_hlsl_invalid_register_or_packoffset);
2042 if ((Component * 32 + Size) > 128) {
2043 Diag(AL.
getLoc(), diag::err_hlsl_packoffset_cross_reg_boundary);
2048 EltTy = VT->getElementType();
2050 if (Align > 32 && Component == 1) {
2053 Diag(AL.
getLoc(), diag::err_hlsl_packoffset_alignment_mismatch)
2067 if (!
SemaRef.checkStringLiteralArgumentAttr(AL, 0, Str, &ArgLoc))
2070 llvm::Triple::EnvironmentType ShaderType;
2071 if (!HLSLShaderAttr::ConvertStrToEnvironmentType(Str, ShaderType)) {
2072 Diag(AL.
getLoc(), diag::warn_attribute_type_not_supported)
2073 << AL << Str << ArgLoc;
2087 assert(AttrList.size() &&
"expected list of resource attributes");
2094 HLSLAttributedResourceType::Attributes ResAttrs;
2096 bool HasResourceClass =
false;
2097 bool HasResourceDimension =
false;
2098 for (
const Attr *A : AttrList) {
2103 case attr::HLSLResourceClass: {
2105 if (HasResourceClass) {
2107 ? diag::warn_duplicate_attribute_exact
2108 : diag::warn_duplicate_attribute)
2112 ResAttrs.ResourceClass = RC;
2113 HasResourceClass =
true;
2116 case attr::HLSLResourceDimension: {
2117 llvm::dxil::ResourceDimension RD =
2119 if (HasResourceDimension) {
2121 ? diag::warn_duplicate_attribute_exact
2122 : diag::warn_duplicate_attribute)
2126 ResAttrs.ResourceDimension = RD;
2127 HasResourceDimension =
true;
2131 if (ResAttrs.IsROV) {
2135 ResAttrs.IsROV =
true;
2137 case attr::HLSLRawBuffer:
2138 if (ResAttrs.RawBuffer) {
2142 ResAttrs.RawBuffer =
true;
2144 case attr::HLSLIsCounter:
2145 if (ResAttrs.IsCounter) {
2149 ResAttrs.IsCounter =
true;
2151 case attr::HLSLContainedType: {
2154 if (!ContainedTy.
isNull()) {
2156 ? diag::warn_duplicate_attribute_exact
2157 : diag::warn_duplicate_attribute)
2166 llvm_unreachable(
"unhandled resource attribute type");
2170 if (!HasResourceClass) {
2171 S.
Diag(AttrList.back()->getRange().getEnd(),
2172 diag::err_hlsl_missing_resource_class);
2177 Wrapped, ContainedTy, ResAttrs);
2179 if (LocInfo && ContainedTyInfo) {
2192 if (!T->isHLSLResourceType()) {
2193 Diag(AL.
getLoc(), diag::err_hlsl_attribute_needs_intangible_type)
2208 AttributeCommonInfo::AS_CXX11, 0, false ,
2213 case ParsedAttr::AT_HLSLResourceClass: {
2215 Diag(AL.
getLoc(), diag::err_attribute_argument_type)
2226 if (!HLSLResourceClassAttr::ConvertStrToResourceClass(Identifier, RC)) {
2227 Diag(ArgLoc, diag::warn_attribute_type_not_supported)
2228 <<
"ResourceClass" << Identifier;
2231 A = HLSLResourceClassAttr::Create(
getASTContext(), RC, ACI);
2235 case ParsedAttr::AT_HLSLResourceDimension: {
2236 StringRef Identifier;
2238 if (!
SemaRef.checkStringLiteralArgumentAttr(AL, 0, Identifier, &ArgLoc))
2242 llvm::dxil::ResourceDimension RD;
2243 if (!HLSLResourceDimensionAttr::ConvertStrToResourceDimension(Identifier,
2245 Diag(ArgLoc, diag::warn_attribute_type_not_supported)
2246 <<
"ResourceDimension" << Identifier;
2249 A = HLSLResourceDimensionAttr::Create(
getASTContext(), RD, ACI);
2253 case ParsedAttr::AT_HLSLROV:
2257 case ParsedAttr::AT_HLSLRawBuffer:
2261 case ParsedAttr::AT_HLSLIsCounter:
2265 case ParsedAttr::AT_HLSLContainedType: {
2267 Diag(AL.
getLoc(), diag::err_attribute_wrong_number_arguments) << AL << 1;
2273 assert(TSI &&
"no type source info for attribute argument");
2275 diag::err_incomplete_type))
2277 A = HLSLContainedTypeAttr::Create(
getASTContext(), TSI, ACI);
2282 llvm_unreachable(
"unhandled HLSL attribute");
2285 HLSLResourcesTypeAttrs.emplace_back(A);
2291 if (!HLSLResourcesTypeAttrs.size())
2297 HLSLResourcesTypeAttrs, QT, &LocInfo)) {
2298 const HLSLAttributedResourceType *RT =
2305 LocsForHLSLAttributedResources.insert(std::pair(RT, LocInfo));
2307 HLSLResourcesTypeAttrs.clear();
2315 auto I = LocsForHLSLAttributedResources.find(RT);
2316 if (I != LocsForHLSLAttributedResources.end()) {
2317 LocInfo = I->second;
2318 LocsForHLSLAttributedResources.erase(I);
2327void SemaHLSL::collectResourceBindingsOnUserRecordDecl(
const VarDecl *VD,
2328 const RecordType *RT) {
2329 const RecordDecl *RD = RT->getDecl()->getDefinitionOrSelf();
2336 "incomplete arrays inside user defined types are not supported");
2345 if (
const HLSLAttributedResourceType *AttrResType =
2346 HLSLAttributedResourceType::findHandleTypeOnResource(Ty)) {
2351 Bindings.addDeclBindingInfo(VD, RC);
2352 }
else if (
const RecordType *RT = dyn_cast<RecordType>(Ty)) {
2358 collectResourceBindingsOnUserRecordDecl(VD, RT);
2370 bool SpecifiedSpace) {
2371 int RegTypeNum =
static_cast<int>(RegType);
2374 if (D->
hasAttr<HLSLGroupSharedAddressSpaceAttr>()) {
2375 S.
Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
2380 if (
HLSLBufferDecl *CBufferOrTBuffer = dyn_cast<HLSLBufferDecl>(D)) {
2381 ResourceClass RC = CBufferOrTBuffer->isCBuffer() ? ResourceClass::CBuffer
2382 : ResourceClass::SRV;
2392 assert(
isa<VarDecl>(D) &&
"D is expected to be VarDecl or HLSLBufferDecl");
2396 if (
const HLSLAttributedResourceType *AttrResType =
2397 HLSLAttributedResourceType::findHandleTypeOnResource(
2414 if (SpecifiedSpace && !DeclaredInCOrTBuffer)
2415 S.
Diag(ArgLoc, diag::err_hlsl_space_on_global_constant);
2420 if (RegType == RegisterType::CBuffer)
2421 S.
Diag(ArgLoc, diag::warn_hlsl_deprecated_register_type_b);
2422 else if (RegType != RegisterType::C)
2423 S.
Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
2427 if (RegType == RegisterType::C)
2428 S.
Diag(ArgLoc, diag::warn_hlsl_register_type_c_packoffset);
2430 S.
Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
2440 S.
Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
2448 bool RegisterTypesDetected[5] = {
false};
2449 RegisterTypesDetected[
static_cast<int>(regType)] =
true;
2452 if (HLSLResourceBindingAttr *
attr =
2453 dyn_cast<HLSLResourceBindingAttr>(*it)) {
2456 if (RegisterTypesDetected[
static_cast<int>(otherRegType)]) {
2457 int otherRegTypeNum =
static_cast<int>(otherRegType);
2459 diag::err_hlsl_duplicate_register_annotation)
2463 RegisterTypesDetected[
static_cast<int>(otherRegType)] =
true;
2471 bool SpecifiedSpace) {
2476 "expecting VarDecl or HLSLBufferDecl");
2488 const uint64_t &Limit,
2491 uint64_t ArrayCount = 1) {
2496 if (StartSlot > Limit)
2500 if (
const auto *AT = dyn_cast<ArrayType>(T)) {
2503 if (
const auto *CAT = dyn_cast<ConstantArrayType>(AT))
2504 Count = CAT->
getSize().getZExtValue();
2508 ArrayCount * Count);
2512 if (
auto ResTy = dyn_cast<HLSLAttributedResourceType>(T)) {
2515 if (ResTy->getAttrs().ResourceClass != ResClass)
2519 uint64_t EndSlot = StartSlot + ArrayCount - 1;
2520 if (EndSlot > Limit)
2524 StartSlot = EndSlot + 1;
2529 if (
const auto *RT = dyn_cast<RecordType>(T)) {
2532 if (
const auto *CXXRD = dyn_cast<CXXRecordDecl>(RD)) {
2535 ResClass, Ctx, ArrayCount))
2542 ResClass, Ctx, ArrayCount))
2556 const uint64_t Limit = UINT32_MAX;
2557 if (SlotNum > Limit)
2562 if (RegTy == RegisterType::C || RegTy == RegisterType::I)
2565 if (
VarDecl *VD = dyn_cast<VarDecl>(TheDecl)) {
2566 uint64_t BaseSlot = SlotNum;
2574 return (BaseSlot > Limit);
2581 return (SlotNum > Limit);
2584 llvm_unreachable(
"unexpected decl type");
2588 if (
VarDecl *VD = dyn_cast<VarDecl>(TheDecl)) {
2590 if (
const auto *IAT = dyn_cast<IncompleteArrayType>(Ty))
2591 Ty = IAT->getElementType();
2593 diag::err_incomplete_type))
2597 StringRef Slot =
"";
2598 StringRef Space =
"";
2602 Diag(AL.
getLoc(), diag::err_attribute_argument_type)
2612 Diag(AL.
getLoc(), diag::err_attribute_argument_type)
2618 SpaceLoc = Loc->
getLoc();
2621 if (Str.starts_with(
"space")) {
2623 SpaceLoc = Loc->
getLoc();
2632 std::optional<unsigned> SlotNum;
2633 unsigned SpaceNum = 0;
2636 if (!Slot.empty()) {
2638 Diag(SlotLoc, diag::err_hlsl_binding_type_invalid) << Slot.substr(0, 1);
2641 if (RegType == RegisterType::I) {
2642 Diag(SlotLoc, diag::warn_hlsl_deprecated_register_type_i);
2645 const StringRef SlotNumStr = Slot.substr(1);
2650 if (SlotNumStr.getAsInteger(10, N)) {
2651 Diag(SlotLoc, diag::err_hlsl_unsupported_register_number);
2659 Diag(SlotLoc, diag::err_hlsl_register_number_too_large);
2668 if (!Space.starts_with(
"space")) {
2669 Diag(SpaceLoc, diag::err_hlsl_expected_space) << Space;
2672 StringRef SpaceNumStr = Space.substr(5);
2673 if (SpaceNumStr.getAsInteger(10, SpaceNum)) {
2674 Diag(SpaceLoc, diag::err_hlsl_expected_space) << Space;
2679 if (SlotNum.has_value())
2684 HLSLResourceBindingAttr *NewAttr =
2685 HLSLResourceBindingAttr::Create(
getASTContext(), Slot, Space, AL);
2687 NewAttr->setBinding(RegType, SlotNum, SpaceNum);
2711 if (
auto *VD = dyn_cast<ValueDecl>(D))
2713 else if (
auto *TD = dyn_cast<TypedefNameDecl>(D))
2714 Ty = TD->getUnderlyingType();
2723 SemaRef.
Diag(Loc, diag::err_hlsl_matrix_layout_non_matrix) << AttrName;
2730 SemaRef.
Diag(Loc, diag::err_hlsl_matrix_layout_non_matrix) << AttrName;
2741 if (
const auto *Existing = D->
getAttr<HLSLMatrixLayoutAttr>()) {
2743 Diag(AL.
getLoc(), diag::err_hlsl_matrix_layout_conflict)
2745 Diag(Existing->getLoc(), diag::note_conflicting_attribute);
2747 Diag(AL.
getLoc(), diag::warn_duplicate_attribute_exact)
2749 Diag(Existing->getLoc(), diag::note_previous_attribute);
2758 Decl *D,
const HLSLMatrixLayoutAttr *
Attr) {
2805 llvm::DenseMap<const FunctionDecl *, unsigned> ScannedDecls;
2809 llvm::Triple::EnvironmentType CurrentShaderEnvironment;
2810 unsigned CurrentShaderStageBit;
2815 bool ReportOnlyShaderStageIssues;
2818 void SetShaderStageContext(llvm::Triple::EnvironmentType ShaderType) {
2819 static_assert(
sizeof(
unsigned) >= 4);
2820 assert(HLSLShaderAttr::isValidShaderType(ShaderType));
2821 assert((
unsigned)(ShaderType - llvm::Triple::Pixel) < 31 &&
2822 "ShaderType is too big for this bitmap");
2825 unsigned bitmapIndex = ShaderType - llvm::Triple::Pixel;
2826 CurrentShaderEnvironment = ShaderType;
2827 CurrentShaderStageBit = (1 << bitmapIndex);
2830 void SetUnknownShaderStageContext() {
2831 CurrentShaderEnvironment = llvm::Triple::UnknownEnvironment;
2832 CurrentShaderStageBit = (1 << 31);
2835 llvm::Triple::EnvironmentType GetCurrentShaderEnvironment()
const {
2836 return CurrentShaderEnvironment;
2839 bool InUnknownShaderStageContext()
const {
2840 return CurrentShaderEnvironment == llvm::Triple::UnknownEnvironment;
2844 void AddToScannedFunctions(
const FunctionDecl *FD) {
2845 unsigned &ScannedStages = ScannedDecls[FD];
2846 ScannedStages |= CurrentShaderStageBit;
2849 unsigned GetScannedStages(
const FunctionDecl *FD) {
return ScannedDecls[FD]; }
2851 bool WasAlreadyScannedInCurrentStage(
const FunctionDecl *FD) {
2852 return WasAlreadyScannedInCurrentStage(GetScannedStages(FD));
2855 bool WasAlreadyScannedInCurrentStage(
unsigned ScannerStages) {
2856 return ScannerStages & CurrentShaderStageBit;
2859 static bool NeverBeenScanned(
unsigned ScannedStages) {
2860 return ScannedStages == 0;
2864 void HandleFunctionOrMethodRef(FunctionDecl *FD, Expr *RefExpr);
2865 void CheckDeclAvailability(NamedDecl *D,
const AvailabilityAttr *AA,
2867 const AvailabilityAttr *FindAvailabilityAttr(
const Decl *D);
2868 bool HasMatchingEnvironmentOrNone(
const AvailabilityAttr *AA);
2871 DiagnoseHLSLAvailability(Sema &SemaRef)
2873 CurrentShaderEnvironment(llvm::Triple::UnknownEnvironment),
2874 CurrentShaderStageBit(0), ReportOnlyShaderStageIssues(
false) {}
2877 void RunOnTranslationUnit(
const TranslationUnitDecl *TU);
2878 void RunOnFunction(
const FunctionDecl *FD);
2880 bool VisitDeclRefExpr(DeclRefExpr *DRE)
override {
2881 FunctionDecl *FD = llvm::dyn_cast<FunctionDecl>(DRE->
getDecl());
2883 HandleFunctionOrMethodRef(FD, DRE);
2887 bool VisitMemberExpr(MemberExpr *ME)
override {
2888 FunctionDecl *FD = llvm::dyn_cast<FunctionDecl>(ME->
getMemberDecl());
2890 HandleFunctionOrMethodRef(FD, ME);
2895void DiagnoseHLSLAvailability::HandleFunctionOrMethodRef(
FunctionDecl *FD,
2898 "expected DeclRefExpr or MemberExpr");
2902 if (FD->
hasBody(FDWithBody)) {
2903 if (!WasAlreadyScannedInCurrentStage(FDWithBody))
2904 DeclsToScan.push_back(FDWithBody);
2909 const AvailabilityAttr *AA = FindAvailabilityAttr(FD);
2911 CheckDeclAvailability(
2915void DiagnoseHLSLAvailability::RunOnTranslationUnit(
2924 DeclContextsToScan.push_back(TU);
2926 while (!DeclContextsToScan.empty()) {
2927 const DeclContext *DC = DeclContextsToScan.pop_back_val();
2928 for (
auto &D : DC->
decls()) {
2935 if (llvm::dyn_cast<NamespaceDecl>(D) || llvm::dyn_cast<ExportDecl>(D)) {
2936 DeclContextsToScan.push_back(llvm::dyn_cast<DeclContext>(D));
2941 const FunctionDecl *FD = llvm::dyn_cast<FunctionDecl>(D);
2946 if (HLSLShaderAttr *ShaderAttr = FD->
getAttr<HLSLShaderAttr>()) {
2947 SetShaderStageContext(ShaderAttr->getType());
2956 for (
const auto *Redecl : FD->
redecls()) {
2957 if (Redecl->isInExportDeclContext()) {
2964 SetUnknownShaderStageContext();
2972void DiagnoseHLSLAvailability::RunOnFunction(
const FunctionDecl *FD) {
2973 assert(DeclsToScan.empty() &&
"DeclsToScan should be empty");
2974 DeclsToScan.push_back(FD);
2976 while (!DeclsToScan.empty()) {
2984 const unsigned ScannedStages = GetScannedStages(FD);
2985 if (WasAlreadyScannedInCurrentStage(ScannedStages))
2988 ReportOnlyShaderStageIssues = !NeverBeenScanned(ScannedStages);
2990 AddToScannedFunctions(FD);
2995bool DiagnoseHLSLAvailability::HasMatchingEnvironmentOrNone(
2996 const AvailabilityAttr *AA) {
3001 llvm::Triple::EnvironmentType CurrentEnv = GetCurrentShaderEnvironment();
3002 if (CurrentEnv == llvm::Triple::UnknownEnvironment)
3005 llvm::Triple::EnvironmentType AttrEnv =
3006 AvailabilityAttr::getEnvironmentType(IIEnvironment->
getName());
3008 return CurrentEnv == AttrEnv;
3011const AvailabilityAttr *
3012DiagnoseHLSLAvailability::FindAvailabilityAttr(
const Decl *D) {
3013 AvailabilityAttr
const *PartialMatch =
nullptr;
3017 for (
const auto *A : D->
attrs()) {
3018 if (
const auto *Avail = dyn_cast<AvailabilityAttr>(A)) {
3019 const AvailabilityAttr *EffectiveAvail = Avail->getEffectiveAttr();
3020 StringRef AttrPlatform = EffectiveAvail->getPlatform()->getName();
3021 StringRef TargetPlatform =
3025 if (AttrPlatform == TargetPlatform) {
3027 if (HasMatchingEnvironmentOrNone(EffectiveAvail))
3029 PartialMatch = Avail;
3033 return PartialMatch;
3038void DiagnoseHLSLAvailability::CheckDeclAvailability(
NamedDecl *D,
3039 const AvailabilityAttr *AA,
3058 if (ReportOnlyShaderStageIssues)
3064 if (InUnknownShaderStageContext())
3069 bool EnvironmentMatches = HasMatchingEnvironmentOrNone(AA);
3070 VersionTuple Introduced = AA->getIntroduced();
3079 llvm::StringRef PlatformName(
3082 llvm::StringRef CurrentEnvStr =
3083 llvm::Triple::getEnvironmentTypeName(GetCurrentShaderEnvironment());
3085 llvm::StringRef AttrEnvStr =
3086 AA->getEnvironment() ? AA->getEnvironment()->getName() :
"";
3087 bool UseEnvironment = !AttrEnvStr.empty();
3089 if (EnvironmentMatches) {
3090 SemaRef.
Diag(
Range.getBegin(), diag::warn_hlsl_availability)
3091 <<
Range << D << PlatformName << Introduced.getAsString()
3092 << UseEnvironment << CurrentEnvStr;
3094 SemaRef.
Diag(
Range.getBegin(), diag::warn_hlsl_availability_unavailable)
3098 SemaRef.
Diag(D->
getLocation(), diag::note_partial_availability_specified_here)
3099 << D << PlatformName << Introduced.getAsString()
3101 << UseEnvironment << AttrEnvStr << CurrentEnvStr;
3108 if (!DefaultCBufferDecls.empty()) {
3111 DefaultCBufferDecls);
3114 SemaRef.getCurLexicalContext()->addDecl(DefaultCBuffer);
3118 for (
const Decl *VD : DefaultCBufferDecls) {
3119 const HLSLResourceBindingAttr *RBA =
3120 VD->
getAttr<HLSLResourceBindingAttr>();
3121 if (RBA && RBA->hasRegisterSlot() &&
3122 RBA->getRegisterType() == HLSLResourceBindingAttr::RegisterType::C) {
3129 SemaRef.Consumer.HandleTopLevelDecl(DG);
3131 diagnoseAvailabilityViolations(TU);
3140 "expected member expr to have resource record type or array of them");
3146 const Expr *NonConstIndexExpr =
nullptr;
3149 if (
const DeclRefExpr *DRE = dyn_cast<DeclRefExpr>(E)) {
3150 if (!NonConstIndexExpr)
3158 diag::err_hlsl_resource_member_array_access_not_constant);
3162 if (
const auto *ASE = dyn_cast<ArraySubscriptExpr>(E)) {
3163 const Expr *IdxExpr = ASE->getIdx();
3165 NonConstIndexExpr = IdxExpr;
3167 }
else if (
const auto *SubME = dyn_cast<MemberExpr>(E)) {
3168 E = SubME->getBase();
3169 }
else if (
const auto *ICE = dyn_cast<ImplicitCastExpr>(E)) {
3170 E = ICE->getSubExpr();
3172 llvm_unreachable(
"unexpected expr type in resource member access");
3181 SemaRef.Context.getCanonicalType(
SemaRef.Context.getAddrSpaceQualType(
3184 SemaRef.Context.getLValueReferenceType(AddrSpaceType));
3187 SemaRef.Context.DeclarationNames.getCXXConversionFunctionName(
3191 [[maybe_unused]]
bool LookupSucceeded =
3192 SemaRef.LookupQualifiedName(ConvR, RD);
3193 assert(LookupSucceeded);
3202std::optional<ExprResult>
3205 const HLSLAttributedResourceType *ResTy =
3206 HLSLAttributedResourceType::findHandleTypeOnResource(
3207 BaseType.getTypePtr());
3209 ResTy->getAttrs().ResourceClass != llvm::dxil::ResourceClass::CBuffer)
3210 return std::nullopt;
3212 QualType TemplateType = ResTy->getContainedType();
3216 assert(NamedConversionDecl &&
3217 "Could not find conversion function for ConstantBuffer.");
3218 auto *ConversionDecl =
3221 return SemaRef.BuildCXXMemberCallExpr(BaseExpr.
get(), NamedConversionDecl,
3233 TI.
getTriple().getEnvironment() != llvm::Triple::EnvironmentType::Library)
3236 DiagnoseHLSLAvailability(
SemaRef).RunOnTranslationUnit(TU);
3243 for (
unsigned I = 1, N = TheCall->
getNumArgs(); I < N; ++I) {
3246 S->
Diag(TheCall->
getBeginLoc(), diag::err_vec_builtin_incompatible_vector)
3271 for (
unsigned I = 0; I < TheCall->
getNumArgs(); ++I) {
3286 if (!BaseType->isFloat32Type())
3287 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3288 << ArgOrdinal << 5 << 0
3298 BaseType = VT->getElementType();
3300 BaseType = MT->getElementType();
3302 if (!BaseType->isHalfType() && !BaseType->isFloat32Type())
3303 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3304 << ArgOrdinal << 5 << 0
3318 if (!BaseType->isDoubleType()) {
3321 return S->
Diag(Loc, diag::err_builtin_requires_double_type)
3322 << ArgOrdinal << PassedType;
3329 unsigned ArgIndex) {
3330 auto *Arg = TheCall->
getArg(ArgIndex);
3332 if (Arg->IgnoreCasts()->isModifiableLvalue(S->
Context, &OrigLoc) ==
3335 S->
Diag(OrigLoc, diag::error_hlsl_inout_lvalue) << Arg << 0;
3345 if (VecTy->getElementType()->isDoubleType())
3346 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3347 << ArgOrdinal << 1 << 0 << 1
3357 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3358 << ArgOrdinal << 5 << 1
3367 if (VecTy->getElementType()->isUnsignedIntegerType())
3370 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3371 << ArgOrdinal << 4 << 3 << 0
3380 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3381 << ArgOrdinal << 5 << 3
3387 unsigned ArgOrdinal,
unsigned Width) {
3390 ArgTy = VTy->getElementType();
3392 uint64_t ElementBitCount =
3394 if (ElementBitCount != Width) {
3396 diag::err_integer_incorrect_bit_count)
3397 << Width << ElementBitCount;
3414 unsigned ArgIndex) {
3423 diag::err_typecheck_expect_scalar_or_vector)
3424 << ArgType << Scalar;
3431 QualType Scalar,
unsigned ArgIndex) {
3442 if (
const auto *VTy = ArgType->getAs<
VectorType>()) {
3455 diag::err_typecheck_expect_scalar_or_vector_or_matrix)
3456 << ArgType << Scalar;
3461 unsigned ArgIndex) {
3466 if (!(ArgType->isScalarType() ||
3467 (VTy && VTy->getElementType()->isScalarType()))) {
3469 diag::err_typecheck_expect_any_scalar_or_vector)
3479 unsigned ArgIndex) {
3481 assert(ArgIndex < TheCall->getNumArgs());
3489 diag::err_typecheck_expect_any_scalar_or_vector)
3514 diag::err_typecheck_call_different_arg_types)
3533 Arg1ScalarTy = VTy->getElementType();
3537 Arg2ScalarTy = VTy->getElementType();
3540 S->
Diag(Arg1->
getBeginLoc(), diag::err_hlsl_builtin_scalar_vector_mismatch)
3541 << 1 << TheCall->
getCallee() << Arg1Ty << Arg2Ty;
3551 if (Arg1Length > 0 && Arg0Length != Arg1Length) {
3553 diag::err_typecheck_vector_lengths_not_equal)
3559 if (Arg2Length > 0 && Arg0Length != Arg2Length) {
3561 diag::err_typecheck_vector_lengths_not_equal)
3573 assert(TheCall->
getNumArgs() > IndexArgIndex &&
"Index argument missing");
3576 unsigned int ActualDim = 1;
3578 ActualDim = VTy->getNumElements();
3579 IndexTy = VTy->getElementType();
3583 diag::err_typecheck_expect_int)
3589 const HLSLAttributedResourceType *ResTy =
3591 assert(ResTy &&
"Resource argument must be a resource");
3592 HLSLAttributedResourceType::Attributes ResAttrs = ResTy->getAttrs();
3594 unsigned int ExpectedDim = 1;
3595 if (ResAttrs.ResourceDimension != llvm::dxil::ResourceDimension::Unknown)
3598 if (ActualDim != ExpectedDim) {
3600 diag::err_hlsl_builtin_resource_coordinate_dimension_mismatch)
3611 llvm::function_ref<
bool(
const HLSLAttributedResourceType *ResType)> Check =
3615 const HLSLAttributedResourceType *ResTy =
3619 diag::err_typecheck_expect_hlsl_resource)
3623 if (Check && Check(ResTy)) {
3625 diag::err_invalid_hlsl_resource_type)
3633 QualType BaseType,
unsigned ExpectedCount,
3635 unsigned PassedCount = 1;
3637 PassedCount = VecTy->getNumElements();
3639 if (PassedCount != ExpectedCount) {
3642 S->
Diag(Loc, diag::err_typecheck_convert_incompatible)
3654 [](
const HLSLAttributedResourceType *ResType) {
3655 return ResType->getAttrs().ResourceDimension ==
3656 llvm::dxil::ResourceDimension::Unknown;
3662 [](
const HLSLAttributedResourceType *ResType) {
3663 return ResType->getAttrs().ResourceClass !=
3664 llvm::hlsl::ResourceClass::Sampler;
3672 unsigned ExpectedDim =
3700 unsigned NextIdx = 3;
3706 diag::err_typecheck_convert_incompatible)
3714 Expr *ComponentArg = TheCall->
getArg(NextIdx);
3718 diag::err_typecheck_convert_incompatible)
3725 std::optional<llvm::APSInt> ComponentOpt =
3728 int64_t ComponentVal = ComponentOpt->getSExtValue();
3729 if (ComponentVal != 0) {
3732 assert(ComponentVal >= 0 && ComponentVal <= 3 &&
3733 "The component is not in the expected range.");
3735 diag::err_hlsl_gathercmp_invalid_component)
3745 const HLSLAttributedResourceType *ResourceTy =
3748 unsigned ExpectedDim =
3757 assert(ResourceTy->hasContainedType() &&
3758 "Expecting a contained type for resource with a dimension "
3760 QualType ReturnType = ResourceTy->getContainedType();
3764 S.
Diag(TheCall->
getBeginLoc(), diag::err_hlsl_samplecmp_requires_float);
3770 ReturnType = VecTy->getElementType();
3783 [](
const HLSLAttributedResourceType *ResType) {
3784 return ResType->getAttrs().ResourceDimension ==
3785 llvm::dxil::ResourceDimension::Unknown;
3793 unsigned ExpectedDim =
3802 EltTy = VTy->getElementType();
3817 TheCall->
setType(ResourceTy->getContainedType());
3822 unsigned MinArgs, MaxArgs;
3850 const HLSLAttributedResourceType *ResourceTy =
3852 unsigned ExpectedDim =
3855 unsigned NextIdx = 3;
3864 diag::err_typecheck_convert_incompatible)
3899 diag::err_typecheck_convert_incompatible)
3905 assert(ResourceTy->hasContainedType() &&
3906 "Expecting a contained type for resource with a dimension "
3908 QualType ReturnType = ResourceTy->getContainedType();
3911 S.
Diag(TheCall->
getBeginLoc(), diag::err_hlsl_samplecmp_requires_float);
3924 switch (BuiltinID) {
3925 case Builtin::BI__builtin_hlsl_adduint64: {
3926 if (
SemaRef.checkArgCount(TheCall, 2))
3940 if (NumElementsArg != 2 && NumElementsArg != 4) {
3942 << 1 << 64 << NumElementsArg * 32;
3956 case Builtin::BI__builtin_hlsl_resource_getpointer: {
3957 if (
SemaRef.checkArgCountRange(TheCall, 1, 2) ||
3964 QualType ContainedTy = ResourceTy->getContainedType();
3965 auto ReturnType =
SemaRef.Context.getAddrSpaceQualType(
3968 ReturnType =
SemaRef.Context.getPointerType(ReturnType);
3973 case Builtin::BI__builtin_hlsl_resource_getpointer_typed: {
3974 if (
SemaRef.checkArgCount(TheCall, 3) ||
3981 "expected pointer type for second argument");
3988 diag::err_invalid_use_of_array_type);
3992 auto ReturnType =
SemaRef.Context.getAddrSpaceQualType(
3995 ReturnType =
SemaRef.Context.getPointerType(ReturnType);
4000 case Builtin::BI__builtin_hlsl_resource_load_with_status: {
4001 if (
SemaRef.checkArgCount(TheCall, 3) ||
4004 SemaRef.getASTContext().UnsignedIntTy) ||
4006 SemaRef.getASTContext().UnsignedIntTy) ||
4012 QualType ReturnType = ResourceTy->getContainedType();
4017 case Builtin::BI__builtin_hlsl_resource_load_with_status_typed: {
4018 if (
SemaRef.checkArgCount(TheCall, 4) ||
4021 SemaRef.getASTContext().UnsignedIntTy) ||
4023 SemaRef.getASTContext().UnsignedIntTy) ||
4029 "expected pointer type for second argument");
4036 diag::err_invalid_use_of_array_type);
4042 case Builtin::BI__builtin_hlsl_resource_load_level:
4044 case Builtin::BI__builtin_hlsl_resource_sample:
4046 case Builtin::BI__builtin_hlsl_resource_sample_bias:
4048 case Builtin::BI__builtin_hlsl_resource_sample_grad:
4050 case Builtin::BI__builtin_hlsl_resource_sample_level:
4052 case Builtin::BI__builtin_hlsl_resource_sample_cmp:
4054 case Builtin::BI__builtin_hlsl_resource_sample_cmp_level_zero:
4056 case Builtin::BI__builtin_hlsl_resource_calculate_lod:
4057 case Builtin::BI__builtin_hlsl_resource_calculate_lod_unclamped:
4059 case Builtin::BI__builtin_hlsl_resource_gather:
4061 case Builtin::BI__builtin_hlsl_resource_gather_cmp:
4063 case Builtin::BI__builtin_hlsl_resource_uninitializedhandle: {
4064 assert(TheCall->
getNumArgs() == 1 &&
"expected 1 arg");
4070 case Builtin::BI__builtin_hlsl_resource_handlefrombinding: {
4071 assert(TheCall->
getNumArgs() == 6 &&
"expected 6 args");
4077 case Builtin::BI__builtin_hlsl_resource_handlefromimplicitbinding: {
4078 assert(TheCall->
getNumArgs() == 6 &&
"expected 6 args");
4084 case Builtin::BI__builtin_hlsl_resource_counterhandlefromimplicitbinding: {
4085 assert(TheCall->
getNumArgs() == 3 &&
"expected 3 args");
4088 auto *MainResType = MainHandleTy->
getAs<HLSLAttributedResourceType>();
4089 auto MainAttrs = MainResType->getAttrs();
4090 assert(!MainAttrs.IsCounter &&
"cannot create a counter from a counter");
4091 MainAttrs.IsCounter =
true;
4093 MainResType->getWrappedType(), MainResType->getContainedType(),
4097 TheCall->
setType(CounterHandleTy);
4100 case Builtin::BI__builtin_hlsl_and:
4101 case Builtin::BI__builtin_hlsl_or: {
4102 if (
SemaRef.checkArgCount(TheCall, 2))
4116 case Builtin::BI__builtin_hlsl_all:
4117 case Builtin::BI__builtin_hlsl_any: {
4118 if (
SemaRef.checkArgCount(TheCall, 1))
4124 case Builtin::BI__builtin_hlsl_asdouble: {
4125 if (
SemaRef.checkArgCount(TheCall, 2))
4129 SemaRef.Context.UnsignedIntTy,
4134 SemaRef.Context.UnsignedIntTy,
4143 case Builtin::BI__builtin_hlsl_elementwise_clamp: {
4144 if (
SemaRef.BuiltinElementwiseTernaryMath(
4150 case Builtin::BI__builtin_hlsl_dot: {
4152 if (
SemaRef.BuiltinVectorToScalarMath(TheCall))
4158 case Builtin::BI__builtin_hlsl_elementwise_firstbithigh:
4159 case Builtin::BI__builtin_hlsl_elementwise_firstbitlow: {
4160 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
4170 EltTy = VecTy->getElementType();
4171 ResTy =
SemaRef.Context.getExtVectorType(ResTy, VecTy->getNumElements());
4184 case Builtin::BI__builtin_hlsl_select: {
4185 if (
SemaRef.checkArgCount(TheCall, 3))
4193 if (VTy && VTy->getElementType()->isBooleanType() &&
4198 case Builtin::BI__builtin_hlsl_elementwise_saturate:
4199 case Builtin::BI__builtin_hlsl_elementwise_rcp: {
4200 if (
SemaRef.checkArgCount(TheCall, 1))
4206 diag::err_builtin_invalid_arg_type)
4209 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
4213 case Builtin::BI__builtin_hlsl_elementwise_degrees:
4214 case Builtin::BI__builtin_hlsl_elementwise_radians:
4215 case Builtin::BI__builtin_hlsl_elementwise_rsqrt:
4216 case Builtin::BI__builtin_hlsl_elementwise_frac:
4217 case Builtin::BI__builtin_hlsl_elementwise_ddx_coarse:
4218 case Builtin::BI__builtin_hlsl_elementwise_ddy_coarse:
4219 case Builtin::BI__builtin_hlsl_elementwise_ddx_fine:
4220 case Builtin::BI__builtin_hlsl_elementwise_ddy_fine: {
4221 if (
SemaRef.checkArgCount(TheCall, 1))
4226 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
4230 case Builtin::BI__builtin_hlsl_elementwise_isinf:
4231 case Builtin::BI__builtin_hlsl_elementwise_isnan: {
4232 if (
SemaRef.checkArgCount(TheCall, 1))
4237 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
4242 case Builtin::BI__builtin_hlsl_lerp: {
4243 if (
SemaRef.checkArgCount(TheCall, 3))
4250 if (
SemaRef.BuiltinElementwiseTernaryMath(TheCall))
4254 case Builtin::BI__builtin_hlsl_mad: {
4255 if (
SemaRef.BuiltinElementwiseTernaryMath(
4261 case Builtin::BI__builtin_hlsl_mul: {
4262 if (
SemaRef.checkArgCount(TheCall, 2))
4271 if (
const auto *VTy = T->getAs<
VectorType>())
4272 return VTy->getElementType();
4274 return MTy->getElementType();
4278 QualType EltTy0 = getElemType(Ty0);
4287 if (IsVec0 && IsMat1) {
4290 }
else if (IsMat0 && IsVec1) {
4294 assert(IsMat0 && IsMat1);
4304 case Builtin::BI__builtin_hlsl_normalize: {
4305 if (
SemaRef.checkArgCount(TheCall, 1))
4316 case Builtin::BI__builtin_elementwise_fma: {
4317 if (
SemaRef.checkArgCount(TheCall, 3) ||
4332 case Builtin::BI__builtin_hlsl_transpose: {
4333 if (
SemaRef.checkArgCount(TheCall, 1))
4342 << 1 << 3 << 0 << 0 << ArgTy;
4347 MatTy->getElementType(), MatTy->getNumColumns(), MatTy->getNumRows());
4351 case Builtin::BI__builtin_hlsl_elementwise_sign: {
4352 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
4360 case Builtin::BI__builtin_hlsl_step: {
4361 if (
SemaRef.checkArgCount(TheCall, 2))
4373 case Builtin::BI__builtin_hlsl_wave_active_all_equal: {
4374 if (
SemaRef.checkArgCount(TheCall, 1))
4388 unsigned NumElts = VecTy->getNumElements();
4398 case Builtin::BI__builtin_hlsl_wave_active_max:
4399 case Builtin::BI__builtin_hlsl_wave_active_min:
4400 case Builtin::BI__builtin_hlsl_wave_active_sum:
4401 case Builtin::BI__builtin_hlsl_wave_active_product: {
4402 if (
SemaRef.checkArgCount(TheCall, 1))
4415 case Builtin::BI__builtin_hlsl_wave_active_bit_or:
4416 case Builtin::BI__builtin_hlsl_wave_active_bit_xor:
4417 case Builtin::BI__builtin_hlsl_wave_active_bit_and: {
4418 if (
SemaRef.checkArgCount(TheCall, 1))
4433 (VTy && VTy->getElementType()->isIntegerType()))) {
4435 diag::err_builtin_invalid_arg_type)
4436 << ArgTyExpr <<
SemaRef.Context.UnsignedIntTy << 1 << 0 << 0;
4446 case Builtin::BI__builtin_elementwise_bitreverse: {
4454 case Builtin::BI__builtin_hlsl_wave_prefix_count_bits: {
4455 if (
SemaRef.checkArgCount(TheCall, 1))
4460 if (!(
ArgType->isScalarType())) {
4462 diag::err_typecheck_expect_any_scalar_or_vector)
4467 if (!(
ArgType->isBooleanType())) {
4469 diag::err_typecheck_expect_any_scalar_or_vector)
4476 case Builtin::BI__builtin_hlsl_wave_read_lane_at: {
4477 if (
SemaRef.checkArgCount(TheCall, 2))
4485 diag::err_typecheck_convert_incompatible)
4486 << ArgTyIndex <<
SemaRef.Context.UnsignedIntTy << 1 << 0 << 0;
4499 case Builtin::BI__builtin_hlsl_wave_get_lane_index: {
4500 if (
SemaRef.checkArgCount(TheCall, 0))
4504 case Builtin::BI__builtin_hlsl_wave_prefix_sum:
4505 case Builtin::BI__builtin_hlsl_wave_prefix_product: {
4506 if (
SemaRef.checkArgCount(TheCall, 1))
4519 case Builtin::BI__builtin_hlsl_quad_read_across_x:
4520 case Builtin::BI__builtin_hlsl_quad_read_across_y: {
4521 if (
SemaRef.checkArgCount(TheCall, 1))
4533 case Builtin::BI__builtin_hlsl_elementwise_splitdouble: {
4534 if (
SemaRef.checkArgCount(TheCall, 3))
4549 case Builtin::BI__builtin_hlsl_elementwise_clip: {
4550 if (
SemaRef.checkArgCount(TheCall, 1))
4557 case Builtin::BI__builtin_elementwise_acos:
4558 case Builtin::BI__builtin_elementwise_asin:
4559 case Builtin::BI__builtin_elementwise_atan:
4560 case Builtin::BI__builtin_elementwise_atan2:
4561 case Builtin::BI__builtin_elementwise_ceil:
4562 case Builtin::BI__builtin_elementwise_cos:
4563 case Builtin::BI__builtin_elementwise_cosh:
4564 case Builtin::BI__builtin_elementwise_exp:
4565 case Builtin::BI__builtin_elementwise_exp2:
4566 case Builtin::BI__builtin_elementwise_exp10:
4567 case Builtin::BI__builtin_elementwise_floor:
4568 case Builtin::BI__builtin_elementwise_fmod:
4569 case Builtin::BI__builtin_elementwise_log:
4570 case Builtin::BI__builtin_elementwise_log2:
4571 case Builtin::BI__builtin_elementwise_log10:
4572 case Builtin::BI__builtin_elementwise_pow:
4573 case Builtin::BI__builtin_elementwise_roundeven:
4574 case Builtin::BI__builtin_elementwise_sin:
4575 case Builtin::BI__builtin_elementwise_sinh:
4576 case Builtin::BI__builtin_elementwise_sqrt:
4577 case Builtin::BI__builtin_elementwise_tan:
4578 case Builtin::BI__builtin_elementwise_tanh:
4579 case Builtin::BI__builtin_elementwise_trunc: {
4585 case Builtin::BI__builtin_hlsl_buffer_update_counter: {
4586 assert(TheCall->
getNumArgs() == 2 &&
"expected 2 args");
4587 auto checkResTy = [](
const HLSLAttributedResourceType *ResTy) ->
bool {
4588 return !(ResTy->getAttrs().ResourceClass == ResourceClass::UAV &&
4589 ResTy->getAttrs().RawBuffer && ResTy->hasContainedType());
4594 std::optional<llvm::APSInt> Offset =
4596 if (!Offset.has_value() ||
std::abs(Offset->getExtValue()) != 1) {
4598 diag::err_hlsl_expect_arg_const_int_one_or_neg_one)
4604 case Builtin::BI__builtin_hlsl_elementwise_f16tof32: {
4605 if (
SemaRef.checkArgCount(TheCall, 1))
4616 ArgTy = VTy->getElementType();
4619 diag::err_builtin_invalid_arg_type)
4628 case Builtin::BI__builtin_hlsl_elementwise_f32tof16: {
4629 if (
SemaRef.checkArgCount(TheCall, 1))
4644 WorkList.push_back(BaseTy);
4645 while (!WorkList.empty()) {
4646 QualType T = WorkList.pop_back_val();
4647 T = T.getCanonicalType().getUnqualifiedType();
4648 if (
const auto *AT = dyn_cast<ConstantArrayType>(T)) {
4656 for (uint64_t Ct = 0; Ct < AT->
getZExtSize(); ++Ct)
4657 llvm::append_range(List, ElementFields);
4662 if (
const auto *VT = dyn_cast<VectorType>(T)) {
4663 List.insert(List.end(), VT->getNumElements(), VT->getElementType());
4666 if (
const auto *MT = dyn_cast<ConstantMatrixType>(T)) {
4667 List.insert(List.end(), MT->getNumElementsFlattened(),
4668 MT->getElementType());
4671 if (
const auto *RD = T->getAsCXXRecordDecl()) {
4672 if (RD->isStandardLayout())
4673 RD = RD->getStandardLayoutBaseWithFields();
4677 if (RD->
isUnion() || !RD->isAggregate()) {
4683 for (
const auto *FD : RD->
fields())
4684 if (!FD->isUnnamedBitField())
4685 FieldTypes.push_back(FD->
getType());
4687 std::reverse(FieldTypes.begin(), FieldTypes.end());
4688 llvm::append_range(WorkList, FieldTypes);
4692 if (!RD->isStandardLayout()) {
4694 for (
const auto &
Base : RD->bases())
4695 FieldTypes.push_back(
Base.getType());
4696 std::reverse(FieldTypes.begin(), FieldTypes.end());
4697 llvm::append_range(WorkList, FieldTypes);
4732 if (
SemaRef.Context.getTypeSize(QT) / 8 > 16)
4738 int ArraySize = VT->getNumElements();
4743 QualType ElTy = VT->getElementType();
4747 if (
SemaRef.Context.getTypeSize(QT) / 8 > 16)
4763 if (
SemaRef.getASTContext().hasSameType(T1, T2))
4772 return llvm::equal(T1Types, T2Types,
4774 return SemaRef.IsLayoutCompatible(LHS, RHS);
4783 bool HadError =
false;
4785 for (
unsigned i = 0, e =
New->getNumParams(); i != e; ++i) {
4793 const auto *NDAttr = NewParam->
getAttr<HLSLParamModifierAttr>();
4794 unsigned NSpellingIdx = (NDAttr ? NDAttr->getSpellingListIndex() : 0);
4795 const auto *ODAttr = OldParam->
getAttr<HLSLParamModifierAttr>();
4796 unsigned OSpellingIdx = (ODAttr ? ODAttr->getSpellingListIndex() : 0);
4798 if (NSpellingIdx != OSpellingIdx) {
4800 diag::err_hlsl_param_qualifier_mismatch)
4801 << NDAttr << NewParam;
4817 if (
SemaRef.getASTContext().hasSameUnqualifiedType(SrcTy, DestTy))
4832 llvm_unreachable(
"HLSL doesn't support pointers.");
4835 llvm_unreachable(
"HLSL doesn't support complex types.");
4837 llvm_unreachable(
"HLSL doesn't support fixed point types.");
4839 llvm_unreachable(
"Should have returned before this");
4849 llvm_unreachable(
"HLSL doesn't support complex types.");
4851 llvm_unreachable(
"HLSL doesn't support fixed point types.");
4856 llvm_unreachable(
"HLSL doesn't support pointers.");
4858 llvm_unreachable(
"Should have returned before this");
4864 llvm_unreachable(
"HLSL doesn't support pointers.");
4867 llvm_unreachable(
"HLSL doesn't support fixed point types.");
4871 llvm_unreachable(
"HLSL doesn't support complex types.");
4874 llvm_unreachable(
"Unhandled scalar cast");
4895 !(SrcMatTy && SrcMatTy->getNumElementsFlattened() == 1))
4901 SrcTy = SrcMatTy->getElementType();
4906 for (
unsigned I = 0, Size = DestTypes.size(); I < Size; ++I) {
4907 if (DestTypes[I]->isUnionType())
4939 if (SrcTypes.size() < DestTypes.size())
4942 unsigned SrcSize = SrcTypes.size();
4943 unsigned DstSize = DestTypes.size();
4945 for (I = 0; I < DstSize && I < SrcSize; I++) {
4946 if (SrcTypes[I]->isUnionType() || DestTypes[I]->isUnionType())
4954 for (; I < SrcSize; I++) {
4955 if (SrcTypes[I]->isUnionType())
4962 assert(Param->hasAttr<HLSLParamModifierAttr>() &&
4963 "We should not get here without a parameter modifier expression");
4964 const auto *
Attr = Param->getAttr<HLSLParamModifierAttr>();
4971 << Arg << (IsInOut ? 1 : 0);
4977 QualType Ty = Param->getType().getNonLValueExprType(Ctx);
4984 << Arg << (IsInOut ? 1 : 0);
4996 SemaRef.PerformCopyInitialization(Entity, Param->getBeginLoc(), ArgOpV);
5002 auto *OpV =
new (Ctx)
5007 Res =
SemaRef.ActOnBinOp(
SemaRef.getCurScope(), Param->getBeginLoc(),
5008 tok::equal, ArgOpV, OpV);
5024 "Pointer and reference types cannot be inout or out parameters");
5025 Ty =
SemaRef.getASTContext().getLValueReferenceType(Ty);
5041 for (
const auto *FD : RD->
fields()) {
5045 assert(RD->getNumBases() <= 1 &&
5046 "HLSL doesn't support multiple inheritance");
5047 return RD->getNumBases()
5052 if (
const auto *AT = dyn_cast<ArrayType>(Ty)) {
5053 if (
const auto *CAT = dyn_cast<ConstantArrayType>(AT))
5065 bool IsVKPushConstant = IsVulkan && VD->
hasAttr<HLSLVkPushConstantAttr>();
5070 !VD->
hasAttr<HLSLVkConstantIdAttr>() && !IsVKPushConstant &&
5076 if (
Decl->getType().hasAddressSpace())
5079 if (
Decl->getType()->isDependentType())
5091 if (
Decl->
hasAttr<HLSLVkExtBuiltinOutputAttr>()) {
5105 llvm::Triple::Vulkan;
5106 if (IsVulkan &&
Decl->
hasAttr<HLSLVkPushConstantAttr>()) {
5107 if (HasDeclaredAPushConstant)
5113 HasDeclaredAPushConstant =
true;
5140class StructBindingContext {
5143 HLSLResourceBindingAttr *RegBindingsAttrs[4];
5144 unsigned RegBindingOffset[4];
5147 static_assert(
static_cast<unsigned>(RegisterType::SRV) == 0 &&
5148 static_cast<unsigned>(RegisterType::UAV) == 1 &&
5149 static_cast<unsigned>(RegisterType::CBuffer) == 2 &&
5150 static_cast<unsigned>(RegisterType::Sampler) == 3,
5151 "unexpected register type values");
5154 HLSLVkBindingAttr *VkBindingAttr;
5155 unsigned VkBindingOffset;
5160 StructBindingContext(
VarDecl *VD) {
5161 for (
unsigned i = 0; i < 4; ++i) {
5162 RegBindingsAttrs[i] =
nullptr;
5163 RegBindingOffset[i] = 0;
5165 VkBindingAttr =
nullptr;
5166 VkBindingOffset = 0;
5172 if (
auto *RBA = dyn_cast<HLSLResourceBindingAttr>(A)) {
5174 unsigned RegTypeIdx =
static_cast<unsigned>(RegType);
5177 RegBindingsAttrs[RegTypeIdx] = RBA;
5182 if (
auto *VBA = dyn_cast<HLSLVkBindingAttr>(A))
5183 VkBindingAttr = VBA;
5190 Attr *createBindingAttr(SemaHLSL &S, ASTContext &AST,
RegisterType RegType,
5191 unsigned Range,
bool HasCounter) {
5192 assert(
static_cast<unsigned>(RegType) < 4 &&
"unexpected register type");
5194 if (VkBindingAttr) {
5195 unsigned Offset = VkBindingOffset;
5196 VkBindingOffset +=
Range;
5197 return HLSLVkBindingAttr::CreateImplicit(
5198 AST, VkBindingAttr->getBinding() + Offset, VkBindingAttr->getSet(),
5199 VkBindingAttr->getRange());
5202 HLSLResourceBindingAttr *RBA =
5203 RegBindingsAttrs[
static_cast<unsigned>(RegType)];
5204 HLSLResourceBindingAttr *NewAttr =
nullptr;
5206 if (RBA && RBA->hasRegisterSlot()) {
5209 unsigned Offset = RegBindingOffset[
static_cast<unsigned>(RegType)];
5210 RegBindingOffset[
static_cast<unsigned>(RegType)] += Range;
5212 unsigned NewSlotNumber = RBA->getSlotNumber() + Offset;
5213 StringRef NewSlotNumberStr =
5215 NewAttr = HLSLResourceBindingAttr::CreateImplicit(
5216 AST, NewSlotNumberStr, RBA->getSpace(), RBA->getRange());
5217 NewAttr->setBinding(RegType, NewSlotNumber, RBA->getSpaceNumber());
5221 NewAttr = HLSLResourceBindingAttr::CreateImplicit(AST,
"",
"0", {});
5222 NewAttr->setBinding(RegType, std::nullopt,
5223 RBA ? RBA->getSpaceNumber() : 0);
5227 NewAttr->setImplicitCounterBindingOrderID(
5236static void createGlobalResourceDeclForStruct(
5238 QualType ResTy, StructBindingContext &BindingCtx) {
5240 "expected resource type or array of resources");
5251 while (
const auto *AT = dyn_cast<ArrayType>(SingleResTy)) {
5252 const auto *CAT = dyn_cast<ConstantArrayType>(AT);
5257 const HLSLAttributedResourceType *ResHandleTy =
5258 HLSLAttributedResourceType::findHandleTypeOnResource(SingleResTy);
5262 Attr *BindingAttr = BindingCtx.createBindingAttr(
5264 ResDecl->
addAttr(BindingAttr);
5265 ResDecl->
addAttr(InternalLinkageAttr::CreateImplicit(AST));
5274 HLSLAssociatedResourceDeclAttr::CreateImplicit(AST, ResDecl));
5281static void handleArrayOfStructWithResources(
5283 EmbeddedResourceNameBuilder &NameBuilder, StructBindingContext &BindingCtx);
5288static void handleStructWithResources(
Sema &S,
VarDecl *ParentVD,
5290 EmbeddedResourceNameBuilder &NameBuilder,
5291 StructBindingContext &BindingCtx) {
5294 assert(RD->
getNumBases() <= 1 &&
"HLSL doesn't support multiple inheritance");
5301 handleStructWithResources(S, ParentVD, BaseRD, NameBuilder, BindingCtx);
5315 createGlobalResourceDeclForStruct(S, ParentVD, FD->
getLocation(), II,
5318 handleStructWithResources(S, ParentVD, RD, NameBuilder, BindingCtx);
5320 }
else if (
const auto *ArrayTy = dyn_cast<ConstantArrayType>(FDTy)) {
5322 "resource arrays should have been already handled");
5323 handleArrayOfStructWithResources(S, ParentVD, ArrayTy, NameBuilder,
5332handleArrayOfStructWithResources(
Sema &S,
VarDecl *ParentVD,
5334 EmbeddedResourceNameBuilder &NameBuilder,
5335 StructBindingContext &BindingCtx) {
5343 if (!SubCAT && !ElementRD)
5346 for (
unsigned I = 0, E = CAT->
getSize().getZExtValue(); I < E; ++I) {
5349 handleStructWithResources(S, ParentVD, ElementRD, NameBuilder,
5352 handleArrayOfStructWithResources(S, ParentVD, SubCAT, NameBuilder,
5365void SemaHLSL::handleGlobalStructOrArrayOfWithResources(
VarDecl *VD) {
5366 EmbeddedResourceNameBuilder NameBuilder(VD->
getName());
5367 StructBindingContext BindingCtx(VD);
5371 "Expected non-resource struct or array type");
5374 handleStructWithResources(
SemaRef, VD, RD, NameBuilder, BindingCtx);
5378 if (
const auto *CAT = dyn_cast<ConstantArrayType>(VDTy)) {
5379 handleArrayOfStructWithResources(
SemaRef, VD, CAT, NameBuilder, BindingCtx);
5387 if (
SemaRef.RequireCompleteType(
5390 diag::err_typecheck_decl_incomplete_type)) {
5404 DefaultCBufferDecls.push_back(VD);
5409 collectResourceBindingsOnVarDecl(VD);
5411 if (VD->
hasAttr<HLSLVkConstantIdAttr>())
5423 processExplicitBindingsOnDecl(VD);
5461 handleGlobalStructOrArrayOfWithResources(VD);
5465 if (VD->
hasAttr<HLSLGroupSharedAddressSpaceAttr>())
5474 "expected resource record type");
5490 const char *CreateMethodName;
5492 CreateMethodName = HasCounter ?
"__createFromBindingWithImplicitCounter"
5493 :
"__createFromBinding";
5495 CreateMethodName = HasCounter
5496 ?
"__createFromImplicitBindingWithImplicitCounter"
5497 :
"__createFromImplicitBinding";
5502 if (!CreateMethod) {
5507 "create method lookup should always succeed for built-in resource "
5516 Args.push_back(RegSlot);
5524 Args.push_back(OrderId);
5530 Args.push_back(Space);
5534 Args.push_back(RangeSize);
5538 Args.push_back(Index);
5540 StringRef VarName = VD->
getName();
5548 Args.push_back(NameCast);
5556 Args.push_back(CounterId);
5579 SemaRef.CheckCompleteVariableDeclaration(VD);
5585 "expected array of resource records");
5606 lookupMethod(
SemaRef, ResourceDecl,
5607 HasCounter ?
"__createFromBindingWithImplicitCounter"
5608 :
"__createFromBinding",
5612 CreateMethod = lookupMethod(
5614 HasCounter ?
"__createFromImplicitBindingWithImplicitCounter"
5615 :
"__createFromImplicitBinding",
5647std::optional<const DeclBindingInfo *> SemaHLSL::inferGlobalBinding(
Expr *E) {
5648 if (
auto *Ternary = dyn_cast<ConditionalOperator>(E)) {
5649 auto TrueInfo = inferGlobalBinding(Ternary->getTrueExpr());
5650 auto FalseInfo = inferGlobalBinding(Ternary->getFalseExpr());
5651 if (!TrueInfo || !FalseInfo)
5652 return std::nullopt;
5653 if (*TrueInfo != *FalseInfo)
5654 return std::nullopt;
5658 if (
auto *ASE = dyn_cast<ArraySubscriptExpr>(E))
5667 if (
const auto *AttrResType =
5668 HLSLAttributedResourceType::findHandleTypeOnResource(Ty)) {
5670 return Bindings.getDeclBindingInfo(VD, RC);
5677void SemaHLSL::trackLocalResource(
VarDecl *VD,
Expr *E) {
5678 std::optional<const DeclBindingInfo *> ExprBinding = inferGlobalBinding(E);
5681 diag::warn_hlsl_assigning_local_resource_is_not_unique)
5686 if (*ExprBinding ==
nullptr)
5689 auto PrevBinding = Assigns.find(VD);
5690 if (PrevBinding == Assigns.end()) {
5692 Assigns.insert({VD, *ExprBinding});
5697 if (*ExprBinding != PrevBinding->second) {
5699 diag::warn_hlsl_assigning_local_resource_is_not_unique)
5701 SemaRef.Diag(VD->getLocation(), diag::note_var_declared_here) << VD;
5712 "expected LHS to be a resource record or array of resource records");
5713 if (Opc != BO_Assign)
5718 while (
auto *ASE = dyn_cast<ArraySubscriptExpr>(E))
5726 SemaRef.Diag(Loc, diag::err_hlsl_assign_to_global_resource) << VD;
5731 trackLocalResource(VD, RHSExpr);
5739void SemaHLSL::collectResourceBindingsOnVarDecl(
VarDecl *VD) {
5741 "expected global variable that contains HLSL resource");
5744 if (
const HLSLBufferDecl *CBufferOrTBuffer = dyn_cast<HLSLBufferDecl>(VD)) {
5745 Bindings.addDeclBindingInfo(VD, CBufferOrTBuffer->isCBuffer()
5746 ? ResourceClass::CBuffer
5747 : ResourceClass::SRV);
5760 if (
const HLSLAttributedResourceType *AttrResType =
5761 HLSLAttributedResourceType::findHandleTypeOnResource(Ty)) {
5762 Bindings.addDeclBindingInfo(VD, AttrResType->getAttrs().ResourceClass);
5767 if (
const RecordType *RT = dyn_cast<RecordType>(Ty))
5768 collectResourceBindingsOnUserRecordDecl(VD, RT);
5774void SemaHLSL::processExplicitBindingsOnDecl(
VarDecl *VD) {
5777 bool HasBinding =
false;
5778 for (Attr *A : VD->
attrs()) {
5781 if (
auto PA = VD->
getAttr<HLSLVkPushConstantAttr>())
5782 Diag(PA->getLoc(), diag::err_hlsl_attr_incompatible) << A << PA;
5785 HLSLResourceBindingAttr *RBA = dyn_cast<HLSLResourceBindingAttr>(A);
5786 if (!RBA || !RBA->hasRegisterSlot())
5791 assert(RT != RegisterType::I &&
"invalid or obsolete register type should "
5792 "never have an attribute created");
5794 if (RT == RegisterType::C) {
5795 if (Bindings.hasBindingInfoForDecl(VD))
5797 diag::warn_hlsl_user_defined_type_missing_member)
5798 <<
static_cast<int>(RT);
5806 if (DeclBindingInfo *BI = Bindings.getDeclBindingInfo(VD, RC)) {
5811 diag::warn_hlsl_user_defined_type_missing_member)
5812 <<
static_cast<int>(RT);
5820class InitListTransformer {
5824 QualType *DstIt =
nullptr;
5825 Expr **ArgIt =
nullptr;
5831 bool castInitializer(Expr *E) {
5832 assert(DstIt &&
"This should always be something!");
5833 if (DstIt == DestTypes.end()) {
5835 ArgExprs.push_back(E);
5840 DstIt = DestTypes.begin();
5843 Ctx, *DstIt,
false);
5848 ArgExprs.push_back(
Init);
5853 bool buildInitializerListImpl(Expr *E) {
5855 if (
auto *
Init = dyn_cast<InitListExpr>(E)) {
5856 for (
auto *SubInit :
Init->inits())
5857 if (!buildInitializerListImpl(SubInit))
5867 return castInitializer(E);
5881 if (
auto *VecTy = Ty->
getAs<VectorType>()) {
5886 for (uint64_t I = 0; I <
Size; ++I) {
5888 SizeTy, SourceLocation());
5894 if (!castInitializer(ElExpr.
get()))
5899 if (
auto *MTy = Ty->
getAs<ConstantMatrixType>()) {
5900 unsigned Rows = MTy->getNumRows();
5901 unsigned Cols = MTy->getNumColumns();
5902 QualType ElemTy = MTy->getElementType();
5904 for (
unsigned R = 0;
R < Rows; ++
R) {
5905 for (
unsigned C = 0;
C < Cols; ++
C) {
5918 if (!castInitializer(ElExpr.
get()))
5926 if (
auto *ArrTy = dyn_cast<ConstantArrayType>(Ty.
getTypePtr())) {
5930 for (uint64_t I = 0; I <
Size; ++I) {
5932 SizeTy, SourceLocation());
5937 if (!buildInitializerListImpl(ElExpr.
get()))
5944 llvm::SmallVector<CXXRecordDecl *> RecordDecls;
5945 RecordDecls.push_back(RD);
5946 while (RecordDecls.back()->getNumBases()) {
5947 CXXRecordDecl *D = RecordDecls.back();
5949 "HLSL doesn't support multiple inheritance");
5950 RecordDecls.push_back(
5953 while (!RecordDecls.empty()) {
5954 CXXRecordDecl *RD = RecordDecls.pop_back_val();
5955 for (
auto *FD : RD->
fields()) {
5956 if (FD->isUnnamedBitField())
5964 if (!buildInitializerListImpl(Res.
get()))
5972 Expr *generateInitListsImpl(QualType Ty) {
5974 assert(ArgIt != ArgExprs.end() &&
"Something is off in iteration!");
5979 llvm::SmallVector<Expr *>
Inits;
5984 if (
auto *ATy = Ty->
getAs<VectorType>()) {
5985 ElTy = ATy->getElementType();
5986 Size = ATy->getNumElements();
5987 }
else if (
auto *CMTy = Ty->
getAs<ConstantMatrixType>()) {
5988 ElTy = CMTy->getElementType();
5989 Size = CMTy->getNumElementsFlattened();
5992 ElTy = VTy->getElementType();
5993 Size = VTy->getZExtSize();
5995 for (uint64_t I = 0; I <
Size; ++I)
5996 Inits.push_back(generateInitListsImpl(ElTy));
5999 llvm::SmallVector<CXXRecordDecl *> RecordDecls;
6000 RecordDecls.push_back(RD);
6001 while (RecordDecls.back()->getNumBases()) {
6002 CXXRecordDecl *D = RecordDecls.back();
6004 "HLSL doesn't support multiple inheritance");
6005 RecordDecls.push_back(
6008 while (!RecordDecls.empty()) {
6009 CXXRecordDecl *RD = RecordDecls.pop_back_val();
6010 for (
auto *FD : RD->
fields())
6011 if (!FD->isUnnamedBitField())
6016 new (Ctx) InitListExpr(Ctx,
Inits.front()->getBeginLoc(),
Inits,
6017 Inits.back()->getEndLoc(),
false);
6018 NewInit->setType(Ty);
6023 llvm::SmallVector<QualType, 16> DestTypes;
6024 llvm::SmallVector<Expr *, 16> ArgExprs;
6025 InitListTransformer(Sema &SemaRef,
const InitializedEntity &Entity)
6026 : S(SemaRef), Ctx(SemaRef.getASTContext()),
6027 Wrap(Entity.
getType()->isIncompleteArrayType()) {
6028 InitTy = Entity.
getType().getNonReferenceType();
6038 DstIt = DestTypes.begin();
6041 bool buildInitializerList(Expr *E) {
return buildInitializerListImpl(E); }
6043 Expr *generateInitLists() {
6044 assert(!ArgExprs.empty() &&
6045 "Call buildInitializerList to generate argument expressions.");
6046 ArgIt = ArgExprs.begin();
6048 return generateInitListsImpl(InitTy);
6049 llvm::SmallVector<Expr *>
Inits;
6050 while (ArgIt != ArgExprs.end())
6051 Inits.push_back(generateInitListsImpl(InitTy));
6054 new (Ctx) InitListExpr(Ctx,
Inits.front()->getBeginLoc(),
Inits,
6055 Inits.back()->getEndLoc(),
false);
6056 llvm::APInt ArySize(64,
Inits.size());
6058 ArraySizeModifier::Normal, 0));
6070 if (
const ArrayType *AT = dyn_cast<ArrayType>(Ty)) {
6077 if (
const auto *RT = Ty->
getAs<RecordType>()) {
6081 if (
const auto *CXXRD = dyn_cast<CXXRecordDecl>(RD)) {
6101 if (
Init->getType()->isScalarType())
6104 InitListTransformer ILT(
SemaRef, Entity);
6106 for (
unsigned I = 0; I <
Init->getNumInits(); ++I) {
6114 Init->setInit(I, E);
6116 if (!ILT.buildInitializerList(E))
6119 size_t ExpectedSize = ILT.DestTypes.size();
6120 size_t ActualSize = ILT.ArgExprs.size();
6121 if (ExpectedSize == 0 && ActualSize == 0)
6128 InitTy =
SemaRef.getASTContext().removeAddrSpaceQualType(InitTy);
6130 SemaRef.Diag(
Init->getBeginLoc(), diag::err_hlsl_incorrect_num_initializers)
6131 << (int)(ExpectedSize < ActualSize) << InitTy
6132 << ExpectedSize << ActualSize;
6142 assert(ExpectedSize > 0 &&
6143 "The expected size of an incomplete array type must be at least 1.");
6145 ((ActualSize + ExpectedSize - 1) / ExpectedSize) * ExpectedSize;
6153 InitTy =
SemaRef.getASTContext().removeAddrSpaceQualType(InitTy);
6154 if (ExpectedSize != ActualSize) {
6155 int TooManyOrFew = ActualSize > ExpectedSize ? 1 : 0;
6156 SemaRef.Diag(
Init->getBeginLoc(), diag::err_hlsl_incorrect_num_initializers)
6157 << TooManyOrFew << InitTy << ExpectedSize << ActualSize;
6164 Init->resizeInits(Ctx, NewInit->getNumInits());
6165 for (
unsigned I = 0; I < NewInit->getNumInits(); ++I)
6166 Init->updateInit(Ctx, I, NewInit->getInit(I));
6174 S.
Diag(OpLoc, diag::err_builtin_matrix_invalid_member)
6184 StringRef AccessorName = CompName->
getName();
6185 assert(!AccessorName.empty() &&
"Matrix Accessor must have a name");
6187 unsigned Rows = MT->getNumRows();
6188 unsigned Cols = MT->getNumColumns();
6189 bool IsZeroBasedAccessor =
false;
6190 unsigned ChunkLen = 0;
6191 if (AccessorName.size() < 2)
6193 "length 4 for zero based: \'_mRC\' or "
6194 "length 3 for one-based: \'_RC\' accessor",
6197 if (AccessorName[0] ==
'_') {
6198 if (AccessorName[1] ==
'm') {
6199 IsZeroBasedAccessor =
true;
6206 S, AccessorName,
"zero based: \'_mRC\' or one-based: \'_RC\' accessor",
6209 if (AccessorName.size() % ChunkLen != 0) {
6210 const llvm::StringRef
Expected = IsZeroBasedAccessor
6211 ?
"zero based: '_mRC' accessor"
6212 :
"one-based: '_RC' accessor";
6217 auto isDigit = [](
char c) {
return c >=
'0' &&
c <=
'9'; };
6218 auto isZeroBasedIndex = [](
unsigned i) {
return i <= 3; };
6219 auto isOneBasedIndex = [](
unsigned i) {
return i >= 1 && i <= 4; };
6221 bool HasRepeated =
false;
6223 unsigned NumComponents = 0;
6224 const char *Begin = AccessorName.data();
6226 for (
unsigned I = 0, E = AccessorName.size(); I < E; I += ChunkLen) {
6227 const char *Chunk = Begin + I;
6228 char RowChar = 0, ColChar = 0;
6229 if (IsZeroBasedAccessor) {
6231 if (Chunk[0] !=
'_' || Chunk[1] !=
'm') {
6232 char Bad = (Chunk[0] !=
'_') ? Chunk[0] : Chunk[1];
6234 S, StringRef(&Bad, 1),
"\'_m\' prefix",
6241 if (Chunk[0] !=
'_')
6243 S, StringRef(&Chunk[0], 1),
"\'_\' prefix",
6250 bool IsDigitsError =
false;
6252 unsigned BadPos = IsZeroBasedAccessor ? 2 : 1;
6256 IsDigitsError =
true;
6260 unsigned BadPos = IsZeroBasedAccessor ? 3 : 2;
6264 IsDigitsError =
true;
6269 unsigned Row = RowChar -
'0';
6270 unsigned Col = ColChar -
'0';
6272 bool HasIndexingError =
false;
6273 if (IsZeroBasedAccessor) {
6275 if (!isZeroBasedIndex(Row)) {
6276 S.
Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
6278 HasIndexingError =
true;
6280 if (!isZeroBasedIndex(Col)) {
6281 S.
Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
6283 HasIndexingError =
true;
6287 if (!isOneBasedIndex(Row)) {
6288 S.
Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
6290 HasIndexingError =
true;
6292 if (!isOneBasedIndex(Col)) {
6293 S.
Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
6295 HasIndexingError =
true;
6302 if (HasIndexingError)
6308 bool HasBoundsError =
false;
6310 Diag(OpLoc, diag::err_hlsl_matrix_index_out_of_bounds)
6312 HasBoundsError =
true;
6315 Diag(OpLoc, diag::err_hlsl_matrix_index_out_of_bounds)
6317 HasBoundsError =
true;
6322 unsigned FlatIndex = Row * Cols + Col;
6323 if (Seen[FlatIndex])
6325 Seen[FlatIndex] =
true;
6328 if (NumComponents == 0 || NumComponents > 4) {
6329 S.
Diag(OpLoc, diag::err_hlsl_matrix_swizzle_invalid_length)
6334 QualType ElemTy = MT->getElementType();
6335 if (NumComponents == 1)
6341 for (Sema::ExtVectorDeclsType::iterator
6345 if ((*I)->getUnderlyingType() == VT)
6356 trackLocalResource(VDecl,
Init);
6358 const HLSLVkConstantIdAttr *ConstIdAttr =
6359 VDecl->
getAttr<HLSLVkConstantIdAttr>();
6366 if (!
Init->isCXX11ConstantExpr(Context, &InitValue)) {
6376 int ConstantID = ConstIdAttr->getId();
6377 llvm::APInt IDVal(Context.getIntWidth(Context.IntTy), ConstantID);
6379 ConstIdAttr->getLocation());
6383 if (
C->getType()->getCanonicalTypeUnqualified() !=
6387 Context.getTrivialTypeSourceInfo(
6388 Init->getType(),
Init->getExprLoc()),
6407 if (!Params || Params->
size() != 1)
6420 if (
auto *TTP = dyn_cast<TemplateTypeParmDecl>(P)) {
6421 if (TTP->hasDefaultArgument()) {
6422 TemplateArgs.
addArgument(TTP->getDefaultArgument());
6425 }
else if (
auto *NTTP = dyn_cast<NonTypeTemplateParmDecl>(P)) {
6426 if (NTTP->hasDefaultArgument()) {
6427 TemplateArgs.
addArgument(NTTP->getDefaultArgument());
6430 }
else if (
auto *TTPD = dyn_cast<TemplateTemplateParmDecl>(P)) {
6431 if (TTPD->hasDefaultArgument()) {
6432 TemplateArgs.
addArgument(TTPD->getDefaultArgument());
6439 return SemaRef.CheckTemplateIdType(
6441 TemplateArgs,
nullptr,
false);
Defines the clang::ASTContext interface.
Defines enum values for all the target-independent builtin functions.
llvm::dxil::ResourceClass ResourceClass
Defines the C++ Decl subclasses, other than those for templates (found in DeclTemplate....
Defines the clang::IdentifierInfo, clang::IdentifierTable, and clang::Selector interfaces.
Forward-declares and imports various common LLVM datatypes that clang wants to use unqualified.
llvm::SmallVector< std::pair< const MemRegion *, SVal >, 4 > Bindings
static bool CheckArgTypeMatches(Sema *S, Expr *Arg, QualType ExpectedType)
static void BuildFlattenedTypeList(QualType BaseTy, llvm::SmallVectorImpl< QualType > &List)
static bool CheckUnsignedIntRepresentation(Sema *S, SourceLocation Loc, int ArgOrdinal, clang::QualType PassedType)
static bool containsIncompleteArrayType(QualType Ty)
static QualType handleIntegerVectorBinOpConversion(Sema &SemaRef, ExprResult &LHS, ExprResult &RHS, QualType LHSType, QualType RHSType, QualType LElTy, QualType RElTy, bool IsCompAssign)
static bool convertToRegisterType(StringRef Slot, RegisterType *RT)
static StringRef createRegisterString(ASTContext &AST, RegisterType RegType, unsigned N)
static bool CheckWaveActive(Sema *S, CallExpr *TheCall)
static void createHostLayoutStructForBuffer(Sema &S, HLSLBufferDecl *BufDecl)
static void castVector(Sema &S, ExprResult &E, QualType &Ty, unsigned Sz)
static QualType ReportMatrixInvalidMember(Sema &S, StringRef Name, StringRef Expected, SourceLocation OpLoc, SourceLocation CompLoc)
static bool CheckBoolSelect(Sema *S, CallExpr *TheCall)
static unsigned calculateLegacyCbufferFieldAlign(const ASTContext &Context, QualType T)
static bool isZeroSizedArray(const ConstantArrayType *CAT)
static bool DiagnoseHLSLRegisterAttribute(Sema &S, SourceLocation &ArgLoc, Decl *D, RegisterType RegType, bool SpecifiedSpace)
static bool hasConstantBufferLayout(QualType QT)
static FieldDecl * createFieldForHostLayoutStruct(Sema &S, const Type *Ty, IdentifierInfo *II, CXXRecordDecl *LayoutStruct)
static bool CheckUnsignedIntVecRepresentation(Sema *S, SourceLocation Loc, int ArgOrdinal, clang::QualType PassedType)
static bool isInvalidConstantBufferLeafElementType(const Type *Ty)
static bool CheckCalculateLodBuiltin(Sema &S, CallExpr *TheCall)
static Builtin::ID getSpecConstBuiltinId(const Type *Type)
static bool CheckFloatingOrIntRepresentation(Sema *S, SourceLocation Loc, int ArgOrdinal, clang::QualType PassedType)
static const Type * createHostLayoutType(Sema &S, const Type *Ty)
static bool CheckAnyScalarOrVector(Sema *S, CallExpr *TheCall, unsigned ArgIndex)
static const HLSLAttributedResourceType * getResourceArrayHandleType(QualType QT)
static IdentifierInfo * getHostLayoutStructName(Sema &S, NamedDecl *BaseDecl, bool MustBeUnique)
static void addImplicitBindingAttrToDecl(Sema &S, Decl *D, RegisterType RT, uint32_t ImplicitBindingOrderID)
static void SetElementTypeAsReturnType(Sema *S, CallExpr *TheCall, QualType ReturnType)
static unsigned calculateLegacyCbufferSize(const ASTContext &Context, QualType T)
static bool CheckLoadLevelBuiltin(Sema &S, CallExpr *TheCall)
static RegisterType getRegisterType(ResourceClass RC)
static bool ValidateRegisterNumber(uint64_t SlotNum, Decl *TheDecl, ASTContext &Ctx, RegisterType RegTy)
static bool isVkPipelineBuiltin(const ASTContext &AstContext, FunctionDecl *FD, HLSLAppliedSemanticAttr *Semantic, bool IsInput)
static bool CheckVectorElementCount(Sema *S, QualType PassedType, QualType BaseType, unsigned ExpectedCount, SourceLocation Loc)
static bool CheckModifiableLValue(Sema *S, CallExpr *TheCall, unsigned ArgIndex)
static QualType castElement(Sema &S, ExprResult &E, QualType Ty)
static char getRegisterTypeChar(RegisterType RT)
static bool CheckNotBoolScalarOrVector(Sema *S, CallExpr *TheCall, unsigned ArgIndex)
static bool isMatrixOrArrayOfMatrix(const ASTContext &Ctx, QualType QT)
static CXXRecordDecl * findRecordDeclInContext(IdentifierInfo *II, DeclContext *DC)
static bool CheckWavePrefix(Sema *S, CallExpr *TheCall)
static bool CheckExpectedBitWidth(Sema *S, CallExpr *TheCall, unsigned ArgOrdinal, unsigned Width)
static LangAS getLangASFromResourceClass(ResourceClass RC)
static bool diagnoseMatrixLayoutOnNonMatrix(Sema &SemaRef, Decl *D, SourceLocation Loc, const IdentifierInfo *AttrName)
static bool CheckVectorSelect(Sema *S, CallExpr *TheCall)
static QualType handleFloatVectorBinOpConversion(Sema &SemaRef, ExprResult &LHS, ExprResult &RHS, QualType LHSType, QualType RHSType, QualType LElTy, QualType RElTy, bool IsCompAssign)
static ResourceClass getResourceClass(RegisterType RT)
static CXXRecordDecl * createHostLayoutStruct(Sema &S, CXXRecordDecl *StructDecl)
static bool CheckScalarOrVector(Sema *S, CallExpr *TheCall, QualType Scalar, unsigned ArgIndex)
static bool CheckSamplingBuiltin(Sema &S, CallExpr *TheCall, SampleKind Kind)
static bool CheckScalarOrVectorOrMatrix(Sema *S, CallExpr *TheCall, QualType Scalar, unsigned ArgIndex)
static bool CheckFloatRepresentation(Sema *S, SourceLocation Loc, int ArgOrdinal, clang::QualType PassedType)
static bool CheckAnyDoubleRepresentation(Sema *S, SourceLocation Loc, int ArgOrdinal, clang::QualType PassedType)
static bool requiresImplicitBufferLayoutStructure(const CXXRecordDecl *RD)
static bool CheckResourceHandle(Sema *S, CallExpr *TheCall, unsigned ArgIndex, llvm::function_ref< bool(const HLSLAttributedResourceType *ResType)> Check=nullptr)
static void validatePackoffset(Sema &S, HLSLBufferDecl *BufDecl)
static bool IsDefaultBufferConstantDecl(const ASTContext &Ctx, VarDecl *VD)
HLSLResourceBindingAttr::RegisterType RegisterType
static CastKind getScalarCastKind(ASTContext &Ctx, QualType DestTy, QualType SrcTy)
static bool CheckGatherBuiltin(Sema &S, CallExpr *TheCall, bool IsCmp)
static bool isValidWaveSizeValue(unsigned Value)
static bool isResourceRecordTypeOrArrayOf(QualType Ty)
static bool AccumulateHLSLResourceSlots(QualType Ty, uint64_t &StartSlot, const uint64_t &Limit, const ResourceClass ResClass, ASTContext &Ctx, uint64_t ArrayCount=1)
static bool CheckNoDoubleVectors(Sema *S, SourceLocation Loc, int ArgOrdinal, clang::QualType PassedType)
static bool ValidateMultipleRegisterAnnotations(Sema &S, Decl *TheDecl, RegisterType regType)
static bool CheckTextureSamplerAndLocation(Sema &S, CallExpr *TheCall)
static bool DiagnoseLocalRegisterBinding(Sema &S, SourceLocation &ArgLoc, Decl *D, RegisterType RegType, bool SpecifiedSpace)
static bool CheckIndexType(Sema *S, CallExpr *TheCall, unsigned IndexArgIndex)
This file declares semantic analysis for HLSL constructs.
Defines the clang::SourceLocation class and associated facilities.
Defines various enumerations that describe declaration and type specifiers.
C Language Family Type Representation.
Defines the clang::TypeLoc interface and its subclasses.
C Language Family Type Representation.
static const TypeInfo & getInfo(unsigned id)
__device__ __2f16 float c
return(__x > > __y)|(__x<<(32 - __y))
APValue - This class implements a discriminated union of [uninitialized] [APSInt] [APFloat],...
virtual bool HandleTopLevelDecl(DeclGroupRef D)
HandleTopLevelDecl - Handle the specified top-level declaration.
Holds long-lived AST nodes (such as types and decls) that can be referred to throughout the semantic ...
unsigned getIntWidth(QualType T) const
int getIntegerTypeOrder(QualType LHS, QualType RHS) const
Return the highest ranked integer type, see C99 6.3.1.8p1.
QualType getPointerType(QualType T) const
Return the uniqued reference to the type for a pointer to the specified type.
const IncompleteArrayType * getAsIncompleteArrayType(QualType T) const
QualType getConstantArrayType(QualType EltTy, const llvm::APInt &ArySize, const Expr *SizeExpr, ArraySizeModifier ASM, unsigned IndexTypeQuals) const
Return the unique reference to the type for a constant array of the specified element type.
QualType getBaseElementType(const ArrayType *VAT) const
Return the innermost element type of an array type.
int getFloatingTypeOrder(QualType LHS, QualType RHS) const
Compare the rank of the two specified floating point types, ignoring the domain of the type (i....
TypeSourceInfo * getTrivialTypeSourceInfo(QualType T, SourceLocation Loc=SourceLocation()) const
Allocate a TypeSourceInfo where all locations have been initialized to a given location,...
QualType getStringLiteralArrayType(QualType EltTy, unsigned Length) const
Return a type for a constant array for a string literal of the specified element type and length.
uint64_t getTypeSize(QualType T) const
Return the size of the specified (complete) type T, in bits.
CharUnits getTypeSizeInChars(QualType T) const
Return the size of the specified (complete) type T, in characters.
CanQualType UnsignedIntTy
QualType getTypedefType(ElaboratedTypeKeyword Keyword, NestedNameSpecifier Qualifier, const TypedefNameDecl *Decl, QualType UnderlyingType=QualType(), std::optional< bool > TypeMatchesDeclOrNone=std::nullopt) const
Return the unique reference to the type for the specified typedef-name decl.
llvm::StringRef backupStr(llvm::StringRef S) const
QualType getSizeType() const
Return the unique type for "size_t" (C99 7.17), defined in <stddef.h>.
QualType getExtVectorType(QualType VectorType, unsigned NumElts) const
Return the unique reference to an extended vector type of the specified element type and size.
const TargetInfo & getTargetInfo() const
QualType getHLSLAttributedResourceType(QualType Wrapped, QualType Contained, const HLSLAttributedResourceType::Attributes &Attrs)
QualType getAddrSpaceQualType(QualType T, LangAS AddressSpace) const
Return the uniqued reference to the type for an address space qualified type with the specified type ...
CanQualType getCanonicalTagType(const TagDecl *TD) const
static bool hasSameUnqualifiedType(QualType T1, QualType T2)
Determine whether the given types are equivalent after cvr-qualifiers have been removed.
QualType getConstantMatrixType(QualType ElementType, unsigned NumRows, unsigned NumColumns) const
Return the unique reference to the matrix type of the specified element type and size.
unsigned getTypeAlign(QualType T) const
Return the ABI-specified alignment of a (complete) type T, in bits.
Represents an array type, per C99 6.7.5.2 - Array Declarators.
QualType getElementType() const
Attr - This represents one attribute.
attr::Kind getKind() const
SourceLocation getLocation() const
SourceLocation getScopeLoc() const
SourceRange getRange() const
const IdentifierInfo * getScopeName() const
SourceLocation getLoc() const
const IdentifierInfo * getAttrName() const
Represents a base class of a C++ class.
QualType getType() const
Retrieves the type of the base class.
Represents a static or instance method of a struct/union/class.
Represents a C++ struct/union/class.
bool isHLSLIntangible() const
Returns true if the class contains HLSL intangible type, either as a field or in base class.
static CXXRecordDecl * Create(const ASTContext &C, TagKind TK, DeclContext *DC, SourceLocation StartLoc, SourceLocation IdLoc, IdentifierInfo *Id, CXXRecordDecl *PrevDecl=nullptr)
void setBases(CXXBaseSpecifier const *const *Bases, unsigned NumBases)
Sets the base classes of this struct or class.
base_class_iterator bases_end()
void completeDefinition() override
Indicates that the definition of this class is now complete.
unsigned getNumBases() const
Retrieves the number of base classes of this class.
base_class_iterator bases_begin()
bool isEmpty() const
Determine whether this is an empty class in the sense of (C++11 [meta.unary.prop]).
CallExpr - Represents a function call (C99 6.5.2.2, C++ [expr.call]).
Expr * getArg(unsigned Arg)
getArg - Return the specified argument.
SourceLocation getBeginLoc() const
static CallExpr * Create(const ASTContext &Ctx, Expr *Fn, ArrayRef< Expr * > Args, QualType Ty, ExprValueKind VK, SourceLocation RParenLoc, FPOptionsOverride FPFeatures, unsigned MinNumArgs=0, ADLCallKind UsesADL=NotADL)
Create a call expression.
FunctionDecl * getDirectCallee()
If the callee is a FunctionDecl, return it. Otherwise return null.
unsigned getNumArgs() const
getNumArgs - Return the number of actual arguments to this call.
static CanQual< Type > CreateUnsafe(QualType Other)
QualType withConst() const
Retrieves a version of this type with const applied.
const T * getTypePtr() const
Retrieve the underlying type pointer, which refers to a canonical type.
QuantityType getQuantity() const
getQuantity - Get the raw integer representation of this quantity.
Represents the canonical version of C arrays with a specified constant size.
bool isZeroSize() const
Return true if the size is zero.
llvm::APInt getSize() const
Return the constant array size as an APInt.
uint64_t getZExtSize() const
Return the size zero-extended as a uint64_t.
Represents a concrete matrix type with constant number of rows and columns.
unsigned getNumColumns() const
Returns the number of columns in the matrix.
static DeclAccessPair make(NamedDecl *D, AccessSpecifier AS)
DeclContext - This is used only as base class of specific decl types that can act as declaration cont...
lookup_result lookup(DeclarationName Name) const
lookup - Find the declarations (if any) with the given Name in this context.
bool isTranslationUnit() const
void addDecl(Decl *D)
Add the declaration D into this context.
decl_range decls() const
decls_begin/decls_end - Iterate over the declarations stored in this context.
DeclContext * getNonTransparentContext()
A reference to a declared variable, function, enum, etc.
static DeclRefExpr * Create(const ASTContext &Context, NestedNameSpecifierLoc QualifierLoc, SourceLocation TemplateKWLoc, ValueDecl *D, bool RefersToEnclosingVariableOrCapture, SourceLocation NameLoc, QualType T, ExprValueKind VK, NamedDecl *FoundD=nullptr, const TemplateArgumentListInfo *TemplateArgs=nullptr, NonOdrUseReason NOUR=NOUR_None)
Decl - This represents one declaration (or definition), e.g.
ASTContext & getASTContext() const LLVM_READONLY
attr_iterator attr_end() const
bool isImplicit() const
isImplicit - Indicates whether the declaration was implicitly generated by the implementation.
void setInvalidDecl(bool Invalid=true)
setInvalidDecl - Indicates the Decl had a semantic error.
bool isInExportDeclContext() const
Whether this declaration was exported in a lexical context.
attr_iterator attr_begin() const
DeclContext * getNonTransparentDeclContext()
Return the non transparent context.
SourceLocation getLocation() const
void setImplicit(bool I=true)
DeclContext * getDeclContext()
AccessSpecifier getAccess() const
SourceLocation getBeginLoc() const LLVM_READONLY
The name of a declaration.
Represents a ValueDecl that came out of a declarator.
SourceLocation getBeginLoc() const LLVM_READONLY
This represents one expression.
bool isIntegerConstantExpr(const ASTContext &Ctx) const
ExprValueKind getValueKind() const
getValueKind - The value kind that this expression produces.
Expr * IgnoreParenImpCasts() LLVM_READONLY
Skip past any parentheses and implicit casts which might surround this expression until reaching a fi...
Expr * IgnoreParens() LLVM_READONLY
Skip past any parentheses which might surround this expression until reaching a fixed point.
std::optional< llvm::APSInt > getIntegerConstantExpr(const ASTContext &Ctx) const
isIntegerConstantExpr - Return the value if this expression is a valid integer constant expression.
bool isLValue() const
isLValue - True if this expression is an "l-value" according to the rules of the current language.
ExprObjectKind getObjectKind() const
getObjectKind - The object kind that this expression produces.
bool HasSideEffects(const ASTContext &Ctx, bool IncludePossibleEffects=true) const
HasSideEffects - This routine returns true for all those expressions which have any effect other than...
SourceLocation getExprLoc() const LLVM_READONLY
getExprLoc - Return the preferred location for the arrow when diagnosing a problem with a generic exp...
ExtVectorType - Extended vector type.
Represents difference between two FPOptions values.
Represents a member of a struct/union/class.
static FieldDecl * Create(const ASTContext &C, DeclContext *DC, SourceLocation StartLoc, SourceLocation IdLoc, const IdentifierInfo *Id, QualType T, TypeSourceInfo *TInfo, Expr *BW, bool Mutable, InClassInitStyle InitStyle)
static FixItHint CreateReplacement(CharSourceRange RemoveRange, StringRef Code)
Create a code modification hint that replaces the given source range with the given code string.
Represents a function declaration or definition.
const ParmVarDecl * getParamDecl(unsigned i) const
Stmt * getBody(const FunctionDecl *&Definition) const
Retrieve the body (definition) of the function.
bool isThisDeclarationADefinition() const
Returns whether this specific declaration of the function is also a definition that does not contain ...
QualType getReturnType() const
ArrayRef< ParmVarDecl * > parameters() const
bool isTemplateInstantiation() const
Determines if the given function was instantiated from a function template.
redecl_range redecls() const
Returns an iterator range for all the redeclarations of the same decl.
unsigned getNumParams() const
Return the number of parameters this function must have based on its FunctionType.
DeclarationNameInfo getNameInfo() const
bool hasBody(const FunctionDecl *&Definition) const
Returns true if the function has a body.
bool isDefined(const FunctionDecl *&Definition, bool CheckForPendingFriendDefinition=false) const
Returns true if the function has a definition that does not need to be instantiated.
Represents a prototype with parameter type info, e.g.
HLSLBufferDecl - Represent a cbuffer or tbuffer declaration.
static HLSLBufferDecl * Create(ASTContext &C, DeclContext *LexicalParent, bool CBuffer, SourceLocation KwLoc, IdentifierInfo *ID, SourceLocation IDLoc, SourceLocation LBrace)
void addLayoutStruct(CXXRecordDecl *LS)
void setHasValidPackoffset(bool PO)
static HLSLBufferDecl * CreateDefaultCBuffer(ASTContext &C, DeclContext *LexicalParent, ArrayRef< Decl * > DefaultCBufferDecls)
buffer_decl_range buffer_decls() const
static HLSLOutArgExpr * Create(const ASTContext &C, QualType Ty, OpaqueValueExpr *Base, OpaqueValueExpr *OpV, Expr *WB, bool IsInOut)
static HLSLRootSignatureDecl * Create(ASTContext &C, DeclContext *DC, SourceLocation Loc, IdentifierInfo *ID, llvm::dxbc::RootSignatureVersion Version, ArrayRef< llvm::hlsl::rootsig::RootElement > RootElements)
One of these records is kept for each identifier that is lexed.
StringRef getName() const
Return the actual identifier string.
A simple pair of identifier info and location.
SourceLocation getLoc() const
IdentifierInfo * getIdentifierInfo() const
IdentifierInfo & get(StringRef Name)
Return the identifier token info for the specified named identifier.
ImplicitCastExpr - Allows us to explicitly represent implicit type conversions, which have no direct ...
static ImplicitCastExpr * Create(const ASTContext &Context, QualType T, CastKind Kind, Expr *Operand, const CXXCastPath *BasePath, ExprValueKind Cat, FPOptionsOverride FPO)
Describes an C or C++ initializer list.
Describes an entity that is being initialized.
QualType getType() const
Retrieve type being initialized.
static InitializedEntity InitializeParameter(ASTContext &Context, ParmVarDecl *Parm)
Create the initialization entity for a parameter.
static IntegerLiteral * Create(const ASTContext &C, const llvm::APInt &V, QualType type, SourceLocation l)
Returns a new integer literal with value 'V' and type 'type'.
iterator begin(Source *source, bool LocalOnly=false)
Represents the results of name lookup.
Represents a prvalue temporary that is written into memory so that a reference can bind to it.
Represents a matrix type, as defined in the Matrix Types clang extensions.
MemberExpr - [C99 6.5.2.3] Structure and Union Members.
ValueDecl * getMemberDecl() const
Retrieve the member declaration to which this expression refers.
This represents a decl that may have a name.
NamedDecl * getUnderlyingDecl()
Looks through UsingDecls and ObjCCompatibleAliasDecls for the underlying named decl.
IdentifierInfo * getIdentifier() const
Get the identifier that names this declaration, if there is one.
StringRef getName() const
Get the name of identifier for this declaration as a StringRef.
DeclarationName getDeclName() const
Get the actual, stored name of the declaration, which may be a special name.
A C++ nested-name-specifier augmented with source location information.
OpaqueValueExpr - An expression referring to an opaque object of a fixed type and value class.
Represents a parameter to a function.
ParsedAttr - Represents a syntactic attribute.
unsigned getSemanticSpelling() const
If the parsed attribute has a semantic equivalent, and it would have a semantic Spelling enumeration ...
unsigned getMinArgs() const
bool checkExactlyNumArgs(class Sema &S, unsigned Num) const
Check if the attribute has exactly as many args as Num.
IdentifierLoc * getArgAsIdent(unsigned Arg) const
bool hasParsedType() const
const ParsedType & getTypeArg() const
unsigned getNumArgs() const
getNumArgs - Return the number of actual arguments to this attribute.
bool isArgIdent(unsigned Arg) const
Expr * getArgAsExpr(unsigned Arg) const
AttributeCommonInfo::Kind getKind() const
A (possibly-)qualified type.
void addRestrict()
Add the restrict qualifier to this QualType.
QualType getNonLValueExprType(const ASTContext &Context) const
Determine the type of a (typically non-lvalue) expression with the specified result type.
QualType getDesugaredType(const ASTContext &Context) const
Return the specified type with any "sugar" removed from the type.
bool isNull() const
Return true if this QualType doesn't point to a type yet.
const Type * getTypePtr() const
Retrieves a pointer to the underlying (unqualified) type.
LangAS getAddressSpace() const
Return the address space of this type.
QualType getNonReferenceType() const
If Type is a reference type (e.g., const int&), returns the type that the reference refers to ("const...
QualType getCanonicalType() const
QualType getUnqualifiedType() const
Retrieve the unqualified variant of the given type, removing as little sugar as possible.
bool hasAddressSpace() const
Check if this type has any address space qualifier.
Represents a struct/union/class.
field_range fields() const
bool hasBindingInfoForDecl(const VarDecl *VD) const
DeclBindingInfo * getDeclBindingInfo(const VarDecl *VD, ResourceClass ResClass)
DeclBindingInfo * addDeclBindingInfo(const VarDecl *VD, ResourceClass ResClass)
Scope - A scope is a transient data structure that is used while parsing the program.
ASTContext & getASTContext() const
SemaDiagnosticBuilder Diag(SourceLocation Loc, unsigned DiagID)
Emit a diagnostic.
ExprResult ActOnOutParamExpr(ParmVarDecl *Param, Expr *Arg)
HLSLRootSignatureDecl * lookupRootSignatureOverrideDecl(DeclContext *DC) const
bool CanPerformElementwiseCast(Expr *Src, QualType DestType)
void handleWaveSizeAttr(Decl *D, const ParsedAttr &AL)
void handleVkLocationAttr(Decl *D, const ParsedAttr &AL)
HLSLAttributedResourceLocInfo TakeLocForHLSLAttribute(const HLSLAttributedResourceType *RT)
void handleSemanticAttr(Decl *D, const ParsedAttr &AL)
bool CanPerformScalarCast(QualType SrcTy, QualType DestTy)
QualType ProcessResourceTypeAttributes(QualType Wrapped)
void handleShaderAttr(Decl *D, const ParsedAttr &AL)
uint32_t getNextImplicitBindingOrderID()
void CheckEntryPoint(FunctionDecl *FD)
void handleVkExtBuiltinOutputAttr(Decl *D, const ParsedAttr &AL)
void emitLogicalOperatorFixIt(Expr *LHS, Expr *RHS, BinaryOperatorKind Opc)
T * createSemanticAttr(const AttributeCommonInfo &ACI, std::optional< unsigned > Location)
bool initGlobalResourceDecl(VarDecl *VD)
void ActOnEndOfTranslationUnit(TranslationUnitDecl *TU)
bool initGlobalResourceArrayDecl(VarDecl *VD)
HLSLVkConstantIdAttr * mergeVkConstantIdAttr(Decl *D, const AttributeCommonInfo &AL, int Id)
HLSLNumThreadsAttr * mergeNumThreadsAttr(Decl *D, const AttributeCommonInfo &AL, int X, int Y, int Z)
void deduceAddressSpace(VarDecl *Decl)
std::pair< IdentifierInfo *, bool > ActOnStartRootSignatureDecl(StringRef Signature)
Computes the unique Root Signature identifier from the given signature, then lookup if there is a pre...
void handlePackOffsetAttr(Decl *D, const ParsedAttr &AL)
std::optional< ExprResult > tryPerformConstantBufferConversion(ExprResult &BaseExpr)
bool diagnosePositionType(QualType T, const ParsedAttr &AL)
bool handleInitialization(VarDecl *VDecl, Expr *&Init)
bool diagnoseInputIDType(QualType T, const ParsedAttr &AL)
void handleParamModifierAttr(Decl *D, const ParsedAttr &AL)
bool CheckResourceBinOp(BinaryOperatorKind Opc, Expr *LHSExpr, Expr *RHSExpr, SourceLocation Loc)
bool CanPerformAggregateSplatCast(Expr *Src, QualType DestType)
bool ActOnResourceMemberAccessExpr(MemberExpr *ME)
bool IsScalarizedLayoutCompatible(QualType T1, QualType T2) const
QualType ActOnTemplateShorthand(TemplateDecl *Template, SourceLocation NameLoc)
void diagnoseSystemSemanticAttr(Decl *D, const ParsedAttr &AL, std::optional< unsigned > Index)
void handleRootSignatureAttr(Decl *D, const ParsedAttr &AL)
bool CheckCompatibleParameterABI(FunctionDecl *New, FunctionDecl *Old)
QualType handleVectorBinOpConversion(ExprResult &LHS, ExprResult &RHS, QualType LHSType, QualType RHSType, bool IsCompAssign)
QualType checkMatrixComponent(Sema &S, QualType baseType, ExprValueKind &VK, SourceLocation OpLoc, const IdentifierInfo *CompName, SourceLocation CompLoc)
bool IsConstantBufferElementCompatible(QualType T1)
void handleResourceBindingAttr(Decl *D, const ParsedAttr &AL)
bool IsTypedResourceElementCompatible(QualType T1)
bool transformInitList(const InitializedEntity &Entity, InitListExpr *Init)
void handleNumThreadsAttr(Decl *D, const ParsedAttr &AL)
bool ActOnUninitializedVarDecl(VarDecl *D)
void handleVkExtBuiltinInputAttr(Decl *D, const ParsedAttr &AL)
bool diagnoseInstantiatedMatrixLayoutAttr(Decl *D, const HLSLMatrixLayoutAttr *Attr)
void ActOnTopLevelFunction(FunctionDecl *FD)
bool handleResourceTypeAttr(QualType T, const ParsedAttr &AL)
void handleVkPushConstantAttr(Decl *D, const ParsedAttr &AL)
HLSLShaderAttr * mergeShaderAttr(Decl *D, const AttributeCommonInfo &AL, llvm::Triple::EnvironmentType ShaderType)
NamedDecl * getConstantBufferConversionFunction(QualType Type, CXXRecordDecl *RD)
void ActOnFinishBuffer(Decl *Dcl, SourceLocation RBrace)
void handleVkBindingAttr(Decl *D, const ParsedAttr &AL)
HLSLParamModifierAttr * mergeParamModifierAttr(Decl *D, const AttributeCommonInfo &AL, HLSLParamModifierAttr::Spelling Spelling)
QualType getInoutParameterType(QualType Ty)
void handleMatrixLayoutAttr(Decl *D, const ParsedAttr &AL)
void handleVkConstantIdAttr(Decl *D, const ParsedAttr &AL)
Decl * ActOnStartBuffer(Scope *BufferScope, bool CBuffer, SourceLocation KwLoc, IdentifierInfo *Ident, SourceLocation IdentLoc, SourceLocation LBrace)
HLSLWaveSizeAttr * mergeWaveSizeAttr(Decl *D, const AttributeCommonInfo &AL, int Min, int Max, int Preferred, int SpelledArgsCount)
bool handleRootSignatureElements(ArrayRef< hlsl::RootSignatureElement > Elements)
void ActOnFinishRootSignatureDecl(SourceLocation Loc, IdentifierInfo *DeclIdent, ArrayRef< hlsl::RootSignatureElement > Elements)
Creates the Root Signature decl of the parsed Root Signature elements onto the AST and push it onto c...
void ActOnVariableDeclarator(VarDecl *VD)
bool CheckBuiltinFunctionCall(unsigned BuiltinID, CallExpr *TheCall)
Sema - This implements semantic analysis and AST building for C.
@ LookupOrdinaryName
Ordinary name lookup, which finds ordinary names (functions, variables, typedefs, etc....
@ LookupMemberName
Member name lookup, which finds the names of class/struct/union members.
ExtVectorDeclsType ExtVectorDecls
ExtVectorDecls - This is a list all the extended vector types.
ASTContext & getASTContext() const
ExprResult ImpCastExprToType(Expr *E, QualType Type, CastKind CK, ExprValueKind VK=VK_PRValue, const CXXCastPath *BasePath=nullptr, CheckedConversionKind CCK=CheckedConversionKind::Implicit)
ImpCastExprToType - If Expr is not of type 'Type', insert an implicit cast.
const LangOptions & getLangOpts() const
ExprResult TemporaryMaterializationConversion(Expr *E)
If E is a prvalue denoting an unmaterialized temporary, materialize it as an xvalue.
ExprResult BuildFieldReferenceExpr(Expr *BaseExpr, bool IsArrow, SourceLocation OpLoc, const CXXScopeSpec &SS, FieldDecl *Field, DeclAccessPair FoundDecl, const DeclarationNameInfo &MemberNameInfo)
bool checkArgCountRange(CallExpr *Call, unsigned MinArgCount, unsigned MaxArgCount)
Checks that a call expression's argument count is in the desired range.
ExternalSemaSource * getExternalSource() const
bool checkArgCount(CallExpr *Call, unsigned DesiredArgCount)
Checks that a call expression's argument count is the desired number.
ExprResult CreateBuiltinArraySubscriptExpr(Expr *Base, SourceLocation LLoc, Expr *Idx, SourceLocation RLoc)
bool LookupQualifiedName(LookupResult &R, DeclContext *LookupCtx, bool InUnqualifiedLookup=false)
Perform qualified name lookup into a given context.
ExprResult PerformCopyInitialization(const InitializedEntity &Entity, SourceLocation EqualLoc, ExprResult Init, bool TopLevelOfInitList=false, bool AllowExplicit=false)
ExprResult CreateBuiltinMatrixSubscriptExpr(Expr *Base, Expr *RowIdx, Expr *ColumnIdx, SourceLocation RBLoc)
Encodes a location in the source.
SourceLocation getLocWithOffset(IntTy Offset) const
Return a source location with the specified offset from this SourceLocation.
A trivial tuple used to represent a source range.
SourceLocation getEnd() const
SourceLocation getEndLoc() const LLVM_READONLY
void printPretty(raw_ostream &OS, PrinterHelper *Helper, const PrintingPolicy &Policy, unsigned Indentation=0, StringRef NewlineSymbol="\n", const ASTContext *Context=nullptr) const
SourceRange getSourceRange() const LLVM_READONLY
SourceLocation tokens are not useful in isolation - they are low level value objects created/interpre...
SourceLocation getBeginLoc() const LLVM_READONLY
StringLiteral - This represents a string literal expression, e.g.
static StringLiteral * Create(const ASTContext &Ctx, StringRef Str, StringLiteralKind Kind, bool Pascal, QualType Ty, ArrayRef< SourceLocation > Locs)
This is the "fully general" constructor that allows representation of strings formed from one or more...
void startDefinition()
Starts the definition of this tag declaration.
Exposes information about the current target.
TargetOptions & getTargetOpts() const
Retrieve the target options.
const llvm::Triple & getTriple() const
Returns the target triple of the primary target.
StringRef getPlatformName() const
Retrieve the name of the platform as it is used in the availability attribute.
VersionTuple getPlatformMinVersion() const
Retrieve the minimum desired version of the platform, to which the program should be compiled.
std::string HLSLEntry
The entry point name for HLSL shader being compiled as specified by -E.
A convenient class for passing around template argument information.
void addArgument(const TemplateArgumentLoc &Loc)
The base class of all kinds of template declarations (e.g., class, function, etc.).
Stores a list of template parameters for a TemplateDecl and its derived classes.
The top declaration context.
SourceLocation getBeginLoc() const
Get the begin source location.
A container of type source information.
TypeLoc getTypeLoc() const
Return the TypeLoc wrapper for the type source info.
The base class of the type hierarchy.
bool isBooleanType() const
bool isIncompleteArrayType() const
CXXRecordDecl * getAsCXXRecordDecl() const
Retrieves the CXXRecordDecl that this type refers to, either because the type is a RecordType or beca...
bool isConstantArrayType() const
bool hasIntegerRepresentation() const
Determine whether this type has an integer representation of some sort, e.g., it is an integer type o...
CXXRecordDecl * castAsCXXRecordDecl() const
bool isArithmeticType() const
bool isConstantMatrixType() const
bool isHLSLBuiltinIntangibleType() const
bool isPointerType() const
CanQualType getCanonicalTypeUnqualified() const
bool isIntegerType() const
isIntegerType() does not include complex integers (a GCC extension).
const T * castAs() const
Member-template castAs<specific type>.
bool isReferenceType() const
bool isHLSLIntangibleType() const
bool isEnumeralType() const
bool isScalarType() const
bool isIntegralType(const ASTContext &Ctx) const
Determine whether this type is an integral type.
const Type * getArrayElementTypeNoTypeQual() const
If this is an array type, return the element type of the array, potentially with type qualifiers miss...
QualType getPointeeType() const
If this is a pointer, ObjC object pointer, or block pointer, this returns the respective pointee.
bool hasUnsignedIntegerRepresentation() const
Determine whether this type has an unsigned integer representation of some sort, e....
bool isDependentType() const
Whether this type is a dependent type, meaning that its definition somehow depends on a template para...
bool isAggregateType() const
Determines whether the type is a C++ aggregate type or C aggregate or union type.
ScalarTypeKind getScalarTypeKind() const
Given that this is a scalar type, classify it.
bool hasSignedIntegerRepresentation() const
Determine whether this type has an signed integer representation of some sort, e.g....
bool isMatrixType() const
bool isHLSLResourceRecord() const
bool hasFloatingRepresentation() const
Determine whether this type has a floating-point representation of some sort, e.g....
bool isVectorType() const
bool isRealFloatingType() const
Floating point categories.
bool isHLSLAttributedResourceType() const
bool isFloatingType() const
const T * getAs() const
Member-template getAs<specific type>'.
const Type * getUnqualifiedDesugaredType() const
Return the specified type with any "sugar" removed from the type, removing any typedefs,...
bool isRecordType() const
bool isHLSLResourceRecordArray() const
void setType(QualType newType)
Represents a variable declaration or definition.
static VarDecl * Create(ASTContext &C, DeclContext *DC, SourceLocation StartLoc, SourceLocation IdLoc, const IdentifierInfo *Id, QualType T, TypeSourceInfo *TInfo, StorageClass S)
void setInitStyle(InitializationStyle Style)
@ CallInit
Call-style initialization (C++98)
void setStorageClass(StorageClass SC)
bool hasGlobalStorage() const
Returns true for all variables that do not have local storage.
StorageClass getStorageClass() const
Returns the storage class as written in the source.
Represents a GCC generic vector type.
unsigned getNumElements() const
QualType getElementType() const
void pushName(llvm::StringRef N)
void pushArrayIndex(uint64_t Index)
void pushBaseName(llvm::StringRef N)
IdentifierInfo * getNameAsIdentifier(ASTContext &AST) const
Defines the clang::TargetInfo interface.
uint32_t getResourceDimensions(llvm::dxil::ResourceDimension Dim)
bool hasCounterHandle(const CXXRecordDecl *RD)
The JSON file list parser is used to communicate input to InstallAPI.
bool isa(CodeGen::Address addr)
if(T->getSizeExpr()) TRY_TO(TraverseStmt(const_cast< Expr * >(T -> getSizeExpr())))
static bool CheckFloatOrHalfRepresentation(Sema *S, SourceLocation Loc, int ArgOrdinal, clang::QualType PassedType)
@ ICIS_NoInit
No in-class initializer.
@ TemplateName
The identifier is a template name. FIXME: Add an annotation for that.
@ OK_Ordinary
An ordinary object is located at an address in memory.
static bool CheckAllArgTypesAreCorrect(Sema *S, CallExpr *TheCall, llvm::ArrayRef< llvm::function_ref< bool(Sema *, SourceLocation, int, QualType)> > Checks)
@ AANT_ArgumentIdentifier
@ Result
The result type of a method or function.
@ Ordinary
This parameter uses ordinary ABI rules for its type.
llvm::Expected< QualType > ExpectedType
@ Template
We are parsing a template declaration.
LLVM_READONLY bool isDigit(unsigned char c)
Return true if this character is an ASCII digit: [0-9].
static bool CheckAllArgsHaveSameType(Sema *S, CallExpr *TheCall)
@ Type
The name was classified as a type.
LangAS
Defines the address space values used by the address space qualifier of QualType.
bool CreateHLSLAttributedResourceType(Sema &S, QualType Wrapped, ArrayRef< const Attr * > AttrList, QualType &ResType, HLSLAttributedResourceLocInfo *LocInfo=nullptr)
CastKind
CastKind - The kind of operation required for a conversion.
ExprValueKind
The categorization of expression values, currently following the C++11 scheme.
@ VK_PRValue
A pr-value expression (in the C++11 taxonomy) produces a temporary value.
@ VK_LValue
An l-value expression is a reference to an object with independent storage.
DynamicRecursiveASTVisitorBase< false > DynamicRecursiveASTVisitor
U cast(CodeGen::Address addr)
@ None
No keyword precedes the qualified type name.
ActionResult< Expr * > ExprResult
Visibility
Describes the different kinds of visibility that a declaration may have.
hash_code hash_value(const clang::dependencies::ModuleID &ID)
__DEVICE__ bool isnan(float __x)
__DEVICE__ _Tp abs(const std::complex< _Tp > &__c)
TypeSourceInfo * ContainedTyInfo
Describes how types, statements, expressions, and declarations should be printed.
unsigned getImplicitOrderID() const
void setCounterImplicitOrderID(unsigned Value) const
bool hasCounterImplicitOrderID() const
unsigned getSpace() const
bool hasImplicitOrderID() const
void setImplicitOrderID(unsigned Value) const
const SourceLocation & getLocation() const
const llvm::hlsl::rootsig::RootElement & getElement() const