37#include "llvm/ADT/ArrayRef.h"
38#include "llvm/ADT/STLExtras.h"
39#include "llvm/ADT/SmallVector.h"
40#include "llvm/ADT/StringExtras.h"
41#include "llvm/ADT/StringRef.h"
42#include "llvm/ADT/Twine.h"
43#include "llvm/Frontend/HLSL/HLSLBinding.h"
44#include "llvm/Frontend/HLSL/RootSignatureValidations.h"
45#include "llvm/Support/Casting.h"
46#include "llvm/Support/DXILABI.h"
47#include "llvm/Support/ErrorHandling.h"
48#include "llvm/Support/FormatVariadic.h"
49#include "llvm/TargetParser/Triple.h"
64 case ResourceClass::SRV:
65 return RegisterType::SRV;
66 case ResourceClass::UAV:
67 return RegisterType::UAV;
68 case ResourceClass::CBuffer:
69 return RegisterType::CBuffer;
70 case ResourceClass::Sampler:
71 return RegisterType::Sampler;
73 llvm_unreachable(
"unexpected ResourceClass value");
83 assert(RT !=
nullptr);
87 *RT = RegisterType::SRV;
91 *RT = RegisterType::UAV;
95 *RT = RegisterType::CBuffer;
99 *RT = RegisterType::Sampler;
103 *RT = RegisterType::C;
107 *RT = RegisterType::I;
116 case RegisterType::SRV:
117 return ResourceClass::SRV;
118 case RegisterType::UAV:
119 return ResourceClass::UAV;
120 case RegisterType::CBuffer:
121 return ResourceClass::CBuffer;
122 case RegisterType::Sampler:
123 return ResourceClass::Sampler;
124 case RegisterType::C:
125 case RegisterType::I:
129 llvm_unreachable(
"unexpected RegisterType value");
133 const auto *BT = dyn_cast<BuiltinType>(
Type);
137 return Builtin::BI__builtin_get_spirv_spec_constant_int;
140 switch (BT->getKind()) {
141 case BuiltinType::Bool:
142 return Builtin::BI__builtin_get_spirv_spec_constant_bool;
143 case BuiltinType::Short:
144 return Builtin::BI__builtin_get_spirv_spec_constant_short;
145 case BuiltinType::Int:
146 return Builtin::BI__builtin_get_spirv_spec_constant_int;
147 case BuiltinType::LongLong:
148 return Builtin::BI__builtin_get_spirv_spec_constant_longlong;
149 case BuiltinType::UShort:
150 return Builtin::BI__builtin_get_spirv_spec_constant_ushort;
151 case BuiltinType::UInt:
152 return Builtin::BI__builtin_get_spirv_spec_constant_uint;
153 case BuiltinType::ULongLong:
154 return Builtin::BI__builtin_get_spirv_spec_constant_ulonglong;
155 case BuiltinType::Half:
156 return Builtin::BI__builtin_get_spirv_spec_constant_half;
157 case BuiltinType::Float:
158 return Builtin::BI__builtin_get_spirv_spec_constant_float;
159 case BuiltinType::Double:
160 return Builtin::BI__builtin_get_spirv_spec_constant_double;
167 ResourceClass ResClass) {
169 "DeclBindingInfo already added");
175 DeclToBindingListIndex.try_emplace(VD, BindingsList.size());
176 return &BindingsList.emplace_back(VD, ResClass);
180 ResourceClass ResClass) {
181 auto Entry = DeclToBindingListIndex.find(VD);
182 if (Entry != DeclToBindingListIndex.end()) {
183 for (
unsigned Index = Entry->getSecond();
184 Index < BindingsList.size() && BindingsList[Index].Decl == VD;
186 if (BindingsList[Index].ResClass == ResClass)
187 return &BindingsList[Index];
194 return DeclToBindingListIndex.contains(VD);
206 getASTContext(), LexicalParent, CBuffer, KwLoc, Ident, IdentLoc, LBrace);
209 auto RC = CBuffer ? llvm::hlsl::ResourceClass::CBuffer
210 : llvm::hlsl::ResourceClass::SRV;
222 if (T->isArrayType() || T->isStructureType() || T->isConstantMatrixType())
229 assert(Context.getTypeSize(T) <= 64 &&
230 "Scalar bit widths larger than 64 not supported");
233 return Context.getTypeSize(T) / 8;
240 constexpr unsigned CBufferAlign = 16;
241 if (
const auto *RD = T->getAsRecordDecl()) {
243 for (
const FieldDecl *Field : RD->fields()) {
250 unsigned AlignSize = llvm::alignTo(Size, FieldAlign);
251 if ((AlignSize % CBufferAlign) + FieldSize > CBufferAlign) {
252 FieldAlign = CBufferAlign;
255 Size = llvm::alignTo(Size, FieldAlign);
262 unsigned ElementCount = AT->getSize().getZExtValue();
263 if (ElementCount == 0)
266 unsigned ElementSize =
268 unsigned AlignedElementSize = llvm::alignTo(ElementSize, CBufferAlign);
269 return AlignedElementSize * (ElementCount - 1) + ElementSize;
273 unsigned ElementCount = VT->getNumElements();
274 unsigned ElementSize =
276 return ElementSize * ElementCount;
279 return Context.getTypeSize(T) / 8;
290 bool HasPackOffset =
false;
291 bool HasNonPackOffset =
false;
293 VarDecl *Var = dyn_cast<VarDecl>(Field);
296 if (Field->hasAttr<HLSLPackOffsetAttr>()) {
297 PackOffsetVec.emplace_back(Var, Field->
getAttr<HLSLPackOffsetAttr>());
298 HasPackOffset =
true;
300 HasNonPackOffset =
true;
307 if (HasNonPackOffset)
314 std::sort(PackOffsetVec.begin(), PackOffsetVec.end(),
315 [](
const std::pair<VarDecl *, HLSLPackOffsetAttr *> &LHS,
316 const std::pair<VarDecl *, HLSLPackOffsetAttr *> &RHS) {
317 return LHS.second->getOffsetInBytes() <
318 RHS.second->getOffsetInBytes();
320 for (
unsigned i = 0; i < PackOffsetVec.size() - 1; i++) {
321 VarDecl *Var = PackOffsetVec[i].first;
322 HLSLPackOffsetAttr *
Attr = PackOffsetVec[i].second;
324 unsigned Begin =
Attr->getOffsetInBytes();
325 unsigned End = Begin + Size;
326 unsigned NextBegin = PackOffsetVec[i + 1].second->getOffsetInBytes();
327 if (End > NextBegin) {
328 VarDecl *NextVar = PackOffsetVec[i + 1].first;
340 CAT = dyn_cast<ConstantArrayType>(
342 return CAT !=
nullptr;
353static const HLSLAttributedResourceType *
356 "expected array of resource records");
358 while (
const ArrayType *AT = dyn_cast<ArrayType>(Ty))
360 return HLSLAttributedResourceType::findHandleTypeOnResource(Ty);
363static const HLSLAttributedResourceType *
377 return RD->isEmpty();
406 Base.getType()->castAsCXXRecordDecl()))
417 assert(RD ==
nullptr &&
418 "there should be at most 1 record by a given name in a scope");
435 Name.append(NameBaseII->
getName());
442 size_t NameLength = Name.size();
451 Name.append(llvm::Twine(suffix).str());
452 II = &AST.
Idents.
get(Name, tok::TokenKind::identifier);
459 Name.truncate(NameLength);
474 if (
const auto *CAT = dyn_cast<ConstantArrayType>(Ty)) {
476 S, CAT->getElementType()->getUnqualifiedDesugaredType());
481 CAT->getSizeModifier(),
482 CAT->getIndexTypeCVRQualifiers())
520 "struct is already HLSL buffer compatible");
534 LS->
addAttr(PackedAttr::CreateImplicit(AST));
538 if (
unsigned NumBases = StructDecl->
getNumBases()) {
539 assert(NumBases == 1 &&
"HLSL supports only one base type");
589 LS->
addAttr(PackedAttr::CreateImplicit(AST));
594 VarDecl *VD = dyn_cast<VarDecl>(D);
610 "host layout field for $Globals decl failed to be created");
627 uint32_t ImplicitBindingOrderID) {
629 HLSLResourceBindingAttr::CreateImplicit(S.
getASTContext(),
"",
"0", {});
630 Attr->setBinding(RT, std::nullopt, 0);
631 Attr->setImplicitBindingOrderID(ImplicitBindingOrderID);
638 BufDecl->setRBraceLoc(RBrace);
655 BufDecl->isCBuffer() ? RegisterType::CBuffer
665 int X,
int Y,
int Z) {
666 if (HLSLNumThreadsAttr *NT = D->
getAttr<HLSLNumThreadsAttr>()) {
667 if (NT->getX() !=
X || NT->getY() != Y || NT->getZ() != Z) {
668 Diag(NT->getLocation(), diag::err_hlsl_attribute_param_mismatch) << AL;
669 Diag(AL.
getLoc(), diag::note_conflicting_attribute);
679 int Min,
int Max,
int Preferred,
680 int SpelledArgsCount) {
681 if (HLSLWaveSizeAttr *WS = D->
getAttr<HLSLWaveSizeAttr>()) {
682 if (WS->getMin() !=
Min || WS->getMax() !=
Max ||
683 WS->getPreferred() != Preferred ||
684 WS->getSpelledArgsCount() != SpelledArgsCount) {
685 Diag(WS->getLocation(), diag::err_hlsl_attribute_param_mismatch) << AL;
686 Diag(AL.
getLoc(), diag::note_conflicting_attribute);
692 Result->setSpelledArgsCount(SpelledArgsCount);
696HLSLVkConstantIdAttr *
702 Diag(AL.
getLoc(), diag::warn_attribute_ignored) << AL;
710 Diag(VD->getLocation(), diag::err_specialization_const);
714 if (!VD->getType().isConstQualified()) {
715 Diag(VD->getLocation(), diag::err_specialization_const);
719 if (HLSLVkConstantIdAttr *CI = D->
getAttr<HLSLVkConstantIdAttr>()) {
720 if (CI->getId() != Id) {
721 Diag(CI->getLocation(), diag::err_hlsl_attribute_param_mismatch) << AL;
722 Diag(AL.
getLoc(), diag::note_conflicting_attribute);
727 HLSLVkConstantIdAttr *
Result =
734 llvm::Triple::EnvironmentType ShaderType) {
735 if (HLSLShaderAttr *NT = D->
getAttr<HLSLShaderAttr>()) {
736 if (NT->getType() != ShaderType) {
737 Diag(NT->getLocation(), diag::err_hlsl_attribute_param_mismatch) << AL;
738 Diag(AL.
getLoc(), diag::note_conflicting_attribute);
742 return HLSLShaderAttr::Create(
getASTContext(), ShaderType, AL);
745HLSLParamModifierAttr *
747 HLSLParamModifierAttr::Spelling Spelling) {
750 if (HLSLParamModifierAttr *PA = D->
getAttr<HLSLParamModifierAttr>()) {
751 if ((PA->isIn() && Spelling == HLSLParamModifierAttr::Keyword_out) ||
752 (PA->isOut() && Spelling == HLSLParamModifierAttr::Keyword_in)) {
753 D->
dropAttr<HLSLParamModifierAttr>();
755 return HLSLParamModifierAttr::Create(
757 HLSLParamModifierAttr::Keyword_inout);
759 Diag(AL.
getLoc(), diag::err_hlsl_duplicate_parameter_modifier) << AL;
760 Diag(PA->getLocation(), diag::note_conflicting_attribute);
786 if (HLSLShaderAttr::isValidShaderType(Env) && Env != llvm::Triple::Library) {
787 if (
const auto *Shader = FD->
getAttr<HLSLShaderAttr>()) {
790 if (Shader->getType() != Env) {
791 Diag(Shader->getLocation(), diag::err_hlsl_entry_shader_attr_mismatch)
803 case llvm::Triple::UnknownEnvironment:
804 case llvm::Triple::Library:
806 case llvm::Triple::RootSignature:
807 llvm_unreachable(
"rootsig environment has no functions");
809 llvm_unreachable(
"Unhandled environment in triple");
815 HLSLAppliedSemanticAttr *Semantic,
820 const auto *ShaderAttr = FD->
getAttr<HLSLShaderAttr>();
821 assert(ShaderAttr &&
"Entry point has no shader attribute");
822 llvm::Triple::EnvironmentType ST = ShaderAttr->getType();
823 auto SemanticName = Semantic->getSemanticName().upper();
828 if (SemanticName ==
"SV_POSITION") {
829 return (ST == llvm::Triple::Vertex && !IsInput) ||
830 (ST == llvm::Triple::Pixel && IsInput);
832 if (SemanticName ==
"SV_VERTEXID")
838bool SemaHLSL::determineActiveSemanticOnScalar(
FunctionDecl *FD,
841 SemanticInfo &ActiveSemantic,
842 SemaHLSL::SemanticContext &SC) {
843 if (ActiveSemantic.Semantic ==
nullptr) {
844 ActiveSemantic.Semantic = D->
getAttr<HLSLParsedSemanticAttr>();
845 if (ActiveSemantic.Semantic)
846 ActiveSemantic.Index = ActiveSemantic.Semantic->getSemanticIndex();
849 if (!ActiveSemantic.Semantic) {
855 HLSLAppliedSemanticAttr(
getASTContext(), *ActiveSemantic.Semantic,
856 ActiveSemantic.Semantic->getAttrName()->getName(),
857 ActiveSemantic.Index.value_or(0));
861 checkSemanticAnnotation(FD, D, A, SC);
862 OutputDecl->addAttr(A);
864 unsigned Location = ActiveSemantic.Index.value_or(0);
867 SC.CurrentIOType & IOType::In)) {
868 bool HasVkLocation =
false;
869 if (
auto *A = D->getAttr<HLSLVkLocationAttr>()) {
870 HasVkLocation = true;
871 Location = A->getLocation();
874 if (SC.UsesExplicitVkLocations.value_or(HasVkLocation) != HasVkLocation) {
875 Diag(D->getLocation(), diag::err_hlsl_semantic_partial_explicit_indexing);
878 SC.UsesExplicitVkLocations = HasVkLocation;
881 const ConstantArrayType *AT = dyn_cast<ConstantArrayType>(D->getType());
882 unsigned ElementCount = AT ? AT->
getZExtSize() : 1;
883 ActiveSemantic.Index = Location + ElementCount;
885 Twine BaseName = Twine(ActiveSemantic.Semantic->getAttrName()->getName());
886 for (
unsigned I = 0; I < ElementCount; ++I) {
887 Twine VariableName = BaseName.concat(Twine(Location + I));
889 auto [_, Inserted] = SC.ActiveSemantics.insert(VariableName.str());
891 Diag(D->getLocation(), diag::err_hlsl_semantic_index_overlap)
892 << VariableName.str();
903 SemanticInfo &ActiveSemantic,
904 SemaHLSL::SemanticContext &SC) {
905 if (ActiveSemantic.Semantic ==
nullptr) {
906 ActiveSemantic.Semantic = D->
getAttr<HLSLParsedSemanticAttr>();
907 if (ActiveSemantic.Semantic)
908 ActiveSemantic.Index = ActiveSemantic.Semantic->getSemanticIndex();
914 const RecordType *RT = dyn_cast<RecordType>(T);
916 return determineActiveSemanticOnScalar(FD, OutputDecl, D, ActiveSemantic,
919 const RecordDecl *RD = RT->getDecl();
920 for (FieldDecl *Field : RD->
fields()) {
921 SemanticInfo Info = ActiveSemantic;
922 if (!determineActiveSemantic(FD, OutputDecl, Field, Info, SC)) {
923 Diag(
Field->getLocation(), diag::note_hlsl_semantic_used_here) <<
Field;
926 if (ActiveSemantic.Semantic)
927 ActiveSemantic = Info;
934 const auto *ShaderAttr = FD->
getAttr<HLSLShaderAttr>();
935 assert(ShaderAttr &&
"Entry point has no shader attribute");
936 llvm::Triple::EnvironmentType ST = ShaderAttr->getType();
940 case llvm::Triple::Pixel:
941 case llvm::Triple::Vertex:
942 case llvm::Triple::Geometry:
943 case llvm::Triple::Hull:
944 case llvm::Triple::Domain:
945 case llvm::Triple::RayGeneration:
946 case llvm::Triple::Intersection:
947 case llvm::Triple::AnyHit:
948 case llvm::Triple::ClosestHit:
949 case llvm::Triple::Miss:
950 case llvm::Triple::Callable:
951 if (
const auto *NT = FD->
getAttr<HLSLNumThreadsAttr>()) {
952 diagnoseAttrStageMismatch(NT, ST,
953 {llvm::Triple::Compute,
954 llvm::Triple::Amplification,
955 llvm::Triple::Mesh});
958 if (
const auto *WS = FD->
getAttr<HLSLWaveSizeAttr>()) {
959 diagnoseAttrStageMismatch(WS, ST,
960 {llvm::Triple::Compute,
961 llvm::Triple::Amplification,
962 llvm::Triple::Mesh});
967 case llvm::Triple::Compute:
968 case llvm::Triple::Amplification:
969 case llvm::Triple::Mesh:
970 if (!FD->
hasAttr<HLSLNumThreadsAttr>()) {
972 << llvm::Triple::getEnvironmentTypeName(ST);
975 if (
const auto *WS = FD->
getAttr<HLSLWaveSizeAttr>()) {
976 if (Ver < VersionTuple(6, 6)) {
977 Diag(WS->getLocation(), diag::err_hlsl_attribute_in_wrong_shader_model)
980 }
else if (WS->getSpelledArgsCount() > 1 && Ver < VersionTuple(6, 8)) {
983 diag::err_hlsl_attribute_number_arguments_insufficient_shader_model)
984 << WS << WS->getSpelledArgsCount() <<
"6.8";
989 case llvm::Triple::RootSignature:
990 llvm_unreachable(
"rootsig environment has no function entry point");
992 llvm_unreachable(
"Unhandled environment in triple");
995 SemaHLSL::SemanticContext InputSC = {};
996 InputSC.CurrentIOType = IOType::In;
999 SemanticInfo ActiveSemantic;
1000 ActiveSemantic.Semantic = Param->getAttr<HLSLParsedSemanticAttr>();
1001 if (ActiveSemantic.Semantic)
1002 ActiveSemantic.Index = ActiveSemantic.Semantic->getSemanticIndex();
1005 if (!determineActiveSemantic(FD, Param, Param, ActiveSemantic, InputSC)) {
1006 Diag(Param->getLocation(), diag::note_previous_decl) << Param;
1011 SemanticInfo ActiveSemantic;
1012 SemaHLSL::SemanticContext OutputSC = {};
1013 OutputSC.CurrentIOType = IOType::Out;
1014 ActiveSemantic.Semantic = FD->
getAttr<HLSLParsedSemanticAttr>();
1015 if (ActiveSemantic.Semantic)
1016 ActiveSemantic.Index = ActiveSemantic.Semantic->getSemanticIndex();
1018 determineActiveSemantic(FD, FD, FD, ActiveSemantic, OutputSC);
1021void SemaHLSL::checkSemanticAnnotation(
1023 const HLSLAppliedSemanticAttr *SemanticAttr,
const SemanticContext &SC) {
1024 auto *ShaderAttr = EntryPoint->
getAttr<HLSLShaderAttr>();
1025 assert(ShaderAttr &&
"Entry point has no shader attribute");
1026 llvm::Triple::EnvironmentType ST = ShaderAttr->getType();
1028 auto SemanticName = SemanticAttr->getSemanticName().upper();
1029 if (SemanticName ==
"SV_DISPATCHTHREADID" ||
1030 SemanticName ==
"SV_GROUPINDEX" || SemanticName ==
"SV_GROUPTHREADID" ||
1031 SemanticName ==
"SV_GROUPID") {
1033 if (ST != llvm::Triple::Compute)
1034 diagnoseSemanticStageMismatch(SemanticAttr, ST, SC.CurrentIOType,
1035 {{llvm::Triple::Compute, IOType::In}});
1037 if (SemanticAttr->getSemanticIndex() != 0) {
1038 std::string PrettyName =
1039 "'" + SemanticAttr->getSemanticName().str() +
"'";
1040 Diag(SemanticAttr->getLoc(),
1041 diag::err_hlsl_semantic_indexing_not_supported)
1047 if (SemanticName ==
"SV_POSITION") {
1050 diagnoseSemanticStageMismatch(SemanticAttr, ST, SC.CurrentIOType,
1051 {{llvm::Triple::Vertex, IOType::InOut},
1052 {llvm::Triple::Pixel, IOType::In}});
1055 if (SemanticName ==
"SV_VERTEXID") {
1056 diagnoseSemanticStageMismatch(SemanticAttr, ST, SC.CurrentIOType,
1057 {{llvm::Triple::Vertex, IOType::In}});
1061 if (SemanticName ==
"SV_TARGET") {
1062 diagnoseSemanticStageMismatch(SemanticAttr, ST, SC.CurrentIOType,
1063 {{llvm::Triple::Pixel, IOType::Out}});
1069 if (SemanticAttr->getAttrName()->getName().starts_with_insensitive(
"SV_"))
1070 llvm_unreachable(
"Unknown SemanticAttr");
1073void SemaHLSL::diagnoseAttrStageMismatch(
1074 const Attr *A, llvm::Triple::EnvironmentType Stage,
1075 std::initializer_list<llvm::Triple::EnvironmentType> AllowedStages) {
1076 SmallVector<StringRef, 8> StageStrings;
1077 llvm::transform(AllowedStages, std::back_inserter(StageStrings),
1078 [](llvm::Triple::EnvironmentType ST) {
1080 HLSLShaderAttr::ConvertEnvironmentTypeToStr(ST));
1082 Diag(A->
getLoc(), diag::err_hlsl_attr_unsupported_in_stage)
1083 << A->
getAttrName() << llvm::Triple::getEnvironmentTypeName(Stage)
1084 << (AllowedStages.size() != 1) << join(StageStrings,
", ");
1087void SemaHLSL::diagnoseSemanticStageMismatch(
1088 const Attr *A, llvm::Triple::EnvironmentType Stage, IOType CurrentIOType,
1089 std::initializer_list<SemanticStageInfo> Allowed) {
1091 for (
auto &Case : Allowed) {
1092 if (Case.Stage != Stage)
1095 if (CurrentIOType & Case.AllowedIOTypesMask)
1098 SmallVector<std::string, 8> ValidCases;
1100 Allowed, std::back_inserter(ValidCases), [](SemanticStageInfo Case) {
1101 SmallVector<std::string, 2> ValidType;
1102 if (Case.AllowedIOTypesMask & IOType::In)
1103 ValidType.push_back(
"input");
1104 if (Case.AllowedIOTypesMask & IOType::Out)
1105 ValidType.push_back(
"output");
1107 HLSLShaderAttr::ConvertEnvironmentTypeToStr(Case.Stage)) +
1108 " " + join(ValidType,
"/");
1110 Diag(A->
getLoc(), diag::err_hlsl_semantic_unsupported_iotype_for_stage)
1111 << A->
getAttrName() << (CurrentIOType & IOType::In ?
"input" :
"output")
1112 << llvm::Triple::getEnvironmentTypeName(Case.Stage)
1113 << join(ValidCases,
", ");
1117 SmallVector<StringRef, 8> StageStrings;
1119 Allowed, std::back_inserter(StageStrings), [](SemanticStageInfo Case) {
1121 HLSLShaderAttr::ConvertEnvironmentTypeToStr(Case.Stage));
1124 Diag(A->
getLoc(), diag::err_hlsl_attr_unsupported_in_stage)
1125 << A->
getAttrName() << llvm::Triple::getEnvironmentTypeName(Stage)
1126 << (Allowed.size() != 1) << join(StageStrings,
", ");
1129template <CastKind Kind>
1132 Ty = VTy->getElementType();
1137template <CastKind Kind>
1149 if (LHSFloat && RHSFloat) {
1177 if (LHSSigned == RHSSigned) {
1178 if (IsCompAssign || IntOrder >= 0)
1186 if (IntOrder != (LHSSigned ? 1 : -1)) {
1187 if (IsCompAssign || RHSSigned)
1195 if (Ctx.getIntWidth(LElTy) != Ctx.getIntWidth(RElTy)) {
1196 if (IsCompAssign || LHSSigned)
1212 QualType ElTy = Ctx.getCorrespondingUnsignedType(LHSSigned ? LElTy : RElTy);
1213 QualType NewTy = Ctx.getExtVectorType(
1223 return CK_FloatingCast;
1225 return CK_IntegralCast;
1227 return CK_IntegralToFloating;
1229 return CK_FloatingToIntegral;
1235 bool IsCompAssign) {
1242 if (!LVecTy && IsCompAssign) {
1244 RHS =
SemaRef.ImpCastExprToType(RHS.
get(), RElTy, CK_HLSLVectorTruncation);
1246 if (Ctx.hasSameUnqualifiedType(LHSType, RHSType))
1248 RHS =
SemaRef.ImpCastExprToType(RHS.
get(), LHSType,
1253 unsigned EndSz = std::numeric_limits<unsigned>::max();
1256 LSz = EndSz = LVecTy->getNumElements();
1259 assert(EndSz != std::numeric_limits<unsigned>::max() &&
1260 "one of the above should have had a value");
1264 if (IsCompAssign && LSz != EndSz) {
1266 diag::err_hlsl_vector_compound_assignment_truncation)
1267 << LHSType << RHSType;
1273 if (!IsCompAssign && LVecTy && LVecTy->getNumElements() > EndSz)
1278 if (!IsCompAssign && !LVecTy)
1282 if (Ctx.hasSameUnqualifiedType(LHSType, RHSType))
1283 return Ctx.getCommonSugaredType(LHSType, RHSType);
1291 LElTy, RElTy, IsCompAssign);
1294 "HLSL Vectors can only contain integer or floating point types");
1296 LElTy, RElTy, IsCompAssign);
1301 assert((Opc == BO_LOr || Opc == BO_LAnd) &&
1302 "Called with non-logical operator");
1304 llvm::raw_svector_ostream OS(Buff);
1306 StringRef NewFnName = Opc == BO_LOr ?
"or" :
"and";
1307 OS << NewFnName <<
"(";
1317std::pair<IdentifierInfo *, bool>
1320 std::string IdStr =
"__hlsl_rootsig_decl_" + std::to_string(Hash);
1327 return {DeclIdent,
Found};
1338 for (
auto &RootSigElement : RootElements)
1339 Elements.push_back(RootSigElement.getElement());
1343 DeclIdent,
SemaRef.getLangOpts().HLSLRootSigVer, Elements);
1345 SignatureDecl->setImplicit();
1351 if (RootSigOverrideIdent) {
1354 if (
SemaRef.LookupQualifiedName(R, DC))
1355 return dyn_cast<HLSLRootSignatureDecl>(R.getFoundDecl());
1363struct PerVisibilityBindingChecker {
1366 std::array<llvm::hlsl::BindingInfoBuilder, 8> Builders;
1370 llvm::dxbc::ShaderVisibility Vis;
1375 PerVisibilityBindingChecker(
SemaHLSL *S) : S(S) {}
1377 void trackBinding(llvm::dxbc::ShaderVisibility
Visibility,
1378 llvm::dxil::ResourceClass RC, uint32_t Space,
1379 uint32_t LowerBound, uint32_t UpperBound,
1380 const hlsl::RootSignatureElement *Elem) {
1382 assert(BuilderIndex < Builders.size() &&
1383 "Not enough builders for visibility type");
1384 Builders[BuilderIndex].trackBinding(RC, Space, LowerBound, UpperBound,
1385 static_cast<const void *
>(Elem));
1387 static_assert(llvm::to_underlying(llvm::dxbc::ShaderVisibility::All) == 0,
1388 "'All' visibility must come first");
1389 if (
Visibility == llvm::dxbc::ShaderVisibility::All)
1390 for (
size_t I = 1, E = Builders.size(); I < E; ++I)
1391 Builders[I].trackBinding(RC, Space, LowerBound, UpperBound,
1392 static_cast<const void *
>(Elem));
1394 ElemInfoMap.push_back({Elem,
Visibility,
false});
1397 ElemInfo &
getInfo(
const hlsl::RootSignatureElement *Elem) {
1398 auto It = llvm::lower_bound(
1400 [](
const auto &LHS,
const auto &RHS) {
return LHS.Elem < RHS; });
1401 assert(It->Elem == Elem &&
"Element not in map");
1405 bool checkOverlap() {
1406 llvm::sort(ElemInfoMap, [](
const auto &LHS,
const auto &RHS) {
1407 return LHS.Elem < RHS.Elem;
1410 bool HadOverlap =
false;
1412 using llvm::hlsl::BindingInfoBuilder;
1413 auto ReportOverlap = [
this,
1414 &HadOverlap](
const BindingInfoBuilder &Builder,
1415 const llvm::hlsl::Binding &Reported) {
1419 static_cast<const hlsl::RootSignatureElement *
>(Reported.Cookie);
1420 const llvm::hlsl::Binding &
Previous = Builder.findOverlapping(Reported);
1421 const auto *PrevElem =
1422 static_cast<const hlsl::RootSignatureElement *
>(
Previous.Cookie);
1424 ElemInfo &Info =
getInfo(Elem);
1429 Info.Diagnosed =
true;
1431 ElemInfo &PrevInfo =
getInfo(PrevElem);
1432 llvm::dxbc::ShaderVisibility CommonVis =
1433 Info.Vis == llvm::dxbc::ShaderVisibility::All ? PrevInfo.Vis
1436 this->S->
Diag(Elem->
getLocation(), diag::err_hlsl_resource_range_overlap)
1437 << llvm::to_underlying(Reported.RC) << Reported.LowerBound
1438 << Reported.isUnbounded() << Reported.UpperBound
1443 this->S->
Diag(PrevElem->getLocation(),
1444 diag::note_hlsl_resource_range_here);
1447 for (BindingInfoBuilder &Builder : Builders)
1448 Builder.calculateBindingInfo(ReportOverlap);
1472 if (
const auto *ResTy =
1473 SecondField->
getType()->
getAs<HLSLAttributedResourceType>()) {
1474 return ResTy->getAttrs().IsCounter;
1482 bool HadError =
false;
1483 auto ReportError = [
this, &HadError](
SourceLocation Loc, uint32_t LowerBound,
1484 uint32_t UpperBound) {
1486 this->
Diag(Loc, diag::err_hlsl_invalid_rootsig_value)
1487 << LowerBound << UpperBound;
1494 this->
Diag(Loc, diag::err_hlsl_invalid_rootsig_value)
1495 << llvm::formatv(
"{0:f}", LowerBound).sstr<6>()
1496 << llvm::formatv(
"{0:f}", UpperBound).sstr<6>();
1499 auto VerifyRegister = [ReportError](
SourceLocation Loc, uint32_t Register) {
1500 if (!llvm::hlsl::rootsig::verifyRegisterValue(Register))
1501 ReportError(Loc, 0, 0xfffffffe);
1504 auto VerifySpace = [ReportError](
SourceLocation Loc, uint32_t Space) {
1505 if (!llvm::hlsl::rootsig::verifyRegisterSpace(Space))
1506 ReportError(Loc, 0, 0xffffffef);
1509 const uint32_t Version =
1510 llvm::to_underlying(
SemaRef.getLangOpts().HLSLRootSigVer);
1511 const uint32_t VersionEnum = Version - 1;
1512 auto ReportFlagError = [
this, &HadError, VersionEnum](
SourceLocation Loc) {
1514 this->
Diag(Loc, diag::err_hlsl_invalid_rootsig_flag)
1521 const llvm::hlsl::rootsig::RootElement &Elem = RootSigElem.
getElement();
1522 if (
const auto *Descriptor =
1523 std::get_if<llvm::hlsl::rootsig::RootDescriptor>(&Elem)) {
1524 VerifyRegister(Loc, Descriptor->Reg.Number);
1525 VerifySpace(Loc, Descriptor->Space);
1527 if (!llvm::hlsl::rootsig::verifyRootDescriptorFlag(Version,
1529 ReportFlagError(Loc);
1530 }
else if (
const auto *Constants =
1531 std::get_if<llvm::hlsl::rootsig::RootConstants>(&Elem)) {
1532 VerifyRegister(Loc, Constants->Reg.Number);
1533 VerifySpace(Loc, Constants->Space);
1534 }
else if (
const auto *Sampler =
1535 std::get_if<llvm::hlsl::rootsig::StaticSampler>(&Elem)) {
1536 VerifyRegister(Loc, Sampler->Reg.Number);
1537 VerifySpace(Loc, Sampler->Space);
1540 "By construction, parseFloatParam can't produce a NaN from a "
1541 "float_literal token");
1543 if (!llvm::hlsl::rootsig::verifyMaxAnisotropy(Sampler->MaxAnisotropy))
1544 ReportError(Loc, 0, 16);
1545 if (!llvm::hlsl::rootsig::verifyMipLODBias(Sampler->MipLODBias))
1546 ReportFloatError(Loc, -16.f, 15.99f);
1547 }
else if (
const auto *Clause =
1548 std::get_if<llvm::hlsl::rootsig::DescriptorTableClause>(
1550 VerifyRegister(Loc, Clause->Reg.Number);
1551 VerifySpace(Loc, Clause->Space);
1553 if (!llvm::hlsl::rootsig::verifyNumDescriptors(Clause->NumDescriptors)) {
1557 ReportError(Loc, 1, 0xfffffffe);
1560 if (!llvm::hlsl::rootsig::verifyDescriptorRangeFlag(Version, Clause->Type,
1562 ReportFlagError(Loc);
1566 PerVisibilityBindingChecker BindingChecker(
this);
1567 SmallVector<std::pair<
const llvm::hlsl::rootsig::DescriptorTableClause *,
1572 const llvm::hlsl::rootsig::RootElement &Elem = RootSigElem.
getElement();
1573 if (
const auto *Descriptor =
1574 std::get_if<llvm::hlsl::rootsig::RootDescriptor>(&Elem)) {
1575 uint32_t LowerBound(Descriptor->Reg.Number);
1576 uint32_t UpperBound(LowerBound);
1578 BindingChecker.trackBinding(
1579 Descriptor->Visibility,
1580 static_cast<llvm::dxil::ResourceClass
>(Descriptor->Type),
1581 Descriptor->Space, LowerBound, UpperBound, &RootSigElem);
1582 }
else if (
const auto *Constants =
1583 std::get_if<llvm::hlsl::rootsig::RootConstants>(&Elem)) {
1584 uint32_t LowerBound(Constants->Reg.Number);
1585 uint32_t UpperBound(LowerBound);
1587 BindingChecker.trackBinding(
1588 Constants->Visibility, llvm::dxil::ResourceClass::CBuffer,
1589 Constants->Space, LowerBound, UpperBound, &RootSigElem);
1590 }
else if (
const auto *Sampler =
1591 std::get_if<llvm::hlsl::rootsig::StaticSampler>(&Elem)) {
1592 uint32_t LowerBound(Sampler->Reg.Number);
1593 uint32_t UpperBound(LowerBound);
1595 BindingChecker.trackBinding(
1596 Sampler->Visibility, llvm::dxil::ResourceClass::Sampler,
1597 Sampler->Space, LowerBound, UpperBound, &RootSigElem);
1598 }
else if (
const auto *Clause =
1599 std::get_if<llvm::hlsl::rootsig::DescriptorTableClause>(
1602 UnboundClauses.emplace_back(Clause, &RootSigElem);
1603 }
else if (
const auto *Table =
1604 std::get_if<llvm::hlsl::rootsig::DescriptorTable>(&Elem)) {
1605 assert(UnboundClauses.size() == Table->NumClauses &&
1606 "Number of unbound elements must match the number of clauses");
1607 bool HasAnySampler =
false;
1608 bool HasAnyNonSampler =
false;
1609 uint64_t Offset = 0;
1610 bool IsPrevUnbound =
false;
1611 for (
const auto &[Clause, ClauseElem] : UnboundClauses) {
1613 if (Clause->Type == llvm::dxil::ResourceClass::Sampler)
1614 HasAnySampler =
true;
1616 HasAnyNonSampler =
true;
1618 if (HasAnySampler && HasAnyNonSampler)
1619 Diag(Loc, diag::err_hlsl_invalid_mixed_resources);
1624 if (Clause->NumDescriptors == 0)
1628 Clause->Offset == llvm::hlsl::rootsig::DescriptorTableOffsetAppend;
1630 Offset = Clause->Offset;
1632 uint64_t RangeBound = llvm::hlsl::rootsig::computeRangeBound(
1633 Offset, Clause->NumDescriptors);
1635 if (IsPrevUnbound && IsAppending)
1636 Diag(Loc, diag::err_hlsl_appending_onto_unbound);
1637 else if (!llvm::hlsl::rootsig::verifyNoOverflowedOffset(RangeBound))
1638 Diag(Loc, diag::err_hlsl_offset_overflow) << Offset << RangeBound;
1641 Offset = RangeBound + 1;
1642 IsPrevUnbound = Clause->NumDescriptors ==
1643 llvm::hlsl::rootsig::NumDescriptorsUnbounded;
1646 uint32_t LowerBound(Clause->Reg.Number);
1647 uint32_t UpperBound = llvm::hlsl::rootsig::computeRangeBound(
1648 LowerBound, Clause->NumDescriptors);
1650 BindingChecker.trackBinding(
1652 static_cast<llvm::dxil::ResourceClass
>(Clause->Type), Clause->Space,
1653 LowerBound, UpperBound, ClauseElem);
1655 UnboundClauses.clear();
1659 return BindingChecker.checkOverlap();
1664 Diag(AL.
getLoc(), diag::err_attribute_wrong_number_arguments) << AL << 1;
1669 if (
auto *RS = D->
getAttr<RootSignatureAttr>()) {
1670 if (RS->getSignatureIdent() != Ident) {
1671 Diag(AL.
getLoc(), diag::err_disallowed_duplicate_attribute) << RS;
1675 Diag(AL.
getLoc(), diag::warn_duplicate_attribute_exact) << RS;
1681 if (
auto *SignatureDecl =
1682 dyn_cast<HLSLRootSignatureDecl>(R.getFoundDecl())) {
1689 llvm::VersionTuple SMVersion =
1694 uint32_t ZMax = 1024;
1695 uint32_t ThreadMax = 1024;
1696 if (IsDXIL && SMVersion.getMajor() <= 4) {
1699 }
else if (IsDXIL && SMVersion.getMajor() == 5) {
1709 diag::err_hlsl_numthreads_argument_oor)
1718 diag::err_hlsl_numthreads_argument_oor)
1727 diag::err_hlsl_numthreads_argument_oor)
1732 if (
X * Y * Z > ThreadMax) {
1733 Diag(AL.
getLoc(), diag::err_hlsl_numthreads_invalid) << ThreadMax;
1750 if (SpelledArgsCount == 0 || SpelledArgsCount > 3)
1758 if (SpelledArgsCount > 1 &&
1762 uint32_t Preferred = 0;
1763 if (SpelledArgsCount > 2 &&
1767 if (SpelledArgsCount > 2) {
1770 diag::err_attribute_power_of_two_in_range)
1771 << AL << llvm::dxil::MinWaveSize << llvm::dxil::MaxWaveSize
1776 if (Preferred < Min || Preferred >
Max) {
1778 diag::err_attribute_power_of_two_in_range)
1779 << AL <<
Min <<
Max << Preferred;
1782 }
else if (SpelledArgsCount > 1) {
1785 diag::err_attribute_power_of_two_in_range)
1786 << AL << llvm::dxil::MinWaveSize << llvm::dxil::MaxWaveSize <<
Max;
1790 Diag(AL.
getLoc(), diag::err_attribute_argument_invalid) << AL << 1;
1793 Diag(AL.
getLoc(), diag::warn_attr_min_eq_max) << AL;
1798 diag::err_attribute_power_of_two_in_range)
1799 << AL << llvm::dxil::MinWaveSize << llvm::dxil::MaxWaveSize <<
Min;
1804 HLSLWaveSizeAttr *NewAttr =
1833 uint32_t Binding = 0;
1857 if (!T->hasUnsignedIntegerRepresentation() ||
1858 (VT && VT->getNumElements() > 3)) {
1859 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type)
1860 << AL <<
"uint/uint2/uint3";
1869 if (!T->hasFloatingRepresentation() || (VT && VT->getNumElements() > 4)) {
1870 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type)
1871 << AL <<
"float/float1/float2/float3/float4";
1879 std::optional<unsigned> Index) {
1883 QualType ValueType = VD->getType();
1884 if (
auto *FD = dyn_cast<FunctionDecl>(D))
1887 bool IsOutput =
false;
1888 if (HLSLParamModifierAttr *MA = D->
getAttr<HLSLParamModifierAttr>()) {
1895 if (SemanticName ==
"SV_DISPATCHTHREADID") {
1898 Diag(AL.
getLoc(), diag::err_hlsl_semantic_output_not_supported) << AL;
1899 if (Index.has_value())
1900 Diag(AL.
getLoc(), diag::err_hlsl_semantic_indexing_not_supported) << AL;
1905 if (SemanticName ==
"SV_GROUPINDEX") {
1907 Diag(AL.
getLoc(), diag::err_hlsl_semantic_output_not_supported) << AL;
1908 if (Index.has_value())
1909 Diag(AL.
getLoc(), diag::err_hlsl_semantic_indexing_not_supported) << AL;
1914 if (SemanticName ==
"SV_GROUPTHREADID") {
1917 Diag(AL.
getLoc(), diag::err_hlsl_semantic_output_not_supported) << AL;
1918 if (Index.has_value())
1919 Diag(AL.
getLoc(), diag::err_hlsl_semantic_indexing_not_supported) << AL;
1924 if (SemanticName ==
"SV_GROUPID") {
1927 Diag(AL.
getLoc(), diag::err_hlsl_semantic_output_not_supported) << AL;
1928 if (Index.has_value())
1929 Diag(AL.
getLoc(), diag::err_hlsl_semantic_indexing_not_supported) << AL;
1934 if (SemanticName ==
"SV_POSITION") {
1935 const auto *VT = ValueType->getAs<
VectorType>();
1936 if (!ValueType->hasFloatingRepresentation() ||
1937 (VT && VT->getNumElements() > 4))
1938 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type)
1939 << AL <<
"float/float1/float2/float3/float4";
1944 if (SemanticName ==
"SV_VERTEXID") {
1945 uint64_t SizeInBits =
SemaRef.Context.getTypeSize(ValueType);
1946 if (!ValueType->isUnsignedIntegerType() || SizeInBits != 32)
1947 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type) << AL <<
"uint";
1952 if (SemanticName ==
"SV_TARGET") {
1953 const auto *VT = ValueType->getAs<
VectorType>();
1954 if (!ValueType->hasFloatingRepresentation() ||
1955 (VT && VT->getNumElements() > 4))
1956 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type)
1957 << AL <<
"float/float1/float2/float3/float4";
1962 Diag(AL.
getLoc(), diag::err_hlsl_unknown_semantic) << AL;
1966 uint32_t IndexValue(0), ExplicitIndex(0);
1969 assert(0 &&
"HLSLUnparsedSemantic is expected to have 2 int arguments.");
1971 assert(IndexValue > 0 ? ExplicitIndex :
true);
1972 std::optional<unsigned> Index =
1973 ExplicitIndex ? std::optional<unsigned>(IndexValue) : std::nullopt;
1983 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_ast_node)
1984 << AL <<
"shader constant in a constant buffer";
1988 uint32_t SubComponent;
1998 bool IsAggregateTy = (T->isArrayType() || T->isStructureType());
2003 if (IsAggregateTy) {
2004 Diag(AL.
getLoc(), diag::err_hlsl_invalid_register_or_packoffset);
2008 if ((Component * 32 + Size) > 128) {
2009 Diag(AL.
getLoc(), diag::err_hlsl_packoffset_cross_reg_boundary);
2014 EltTy = VT->getElementType();
2016 if (Align > 32 && Component == 1) {
2019 Diag(AL.
getLoc(), diag::err_hlsl_packoffset_alignment_mismatch)
2033 if (!
SemaRef.checkStringLiteralArgumentAttr(AL, 0, Str, &ArgLoc))
2036 llvm::Triple::EnvironmentType ShaderType;
2037 if (!HLSLShaderAttr::ConvertStrToEnvironmentType(Str, ShaderType)) {
2038 Diag(AL.
getLoc(), diag::warn_attribute_type_not_supported)
2039 << AL << Str << ArgLoc;
2053 assert(AttrList.size() &&
"expected list of resource attributes");
2060 HLSLAttributedResourceType::Attributes ResAttrs;
2062 bool HasResourceClass =
false;
2063 bool HasResourceDimension =
false;
2064 for (
const Attr *A : AttrList) {
2069 case attr::HLSLResourceClass: {
2071 if (HasResourceClass) {
2073 ? diag::warn_duplicate_attribute_exact
2074 : diag::warn_duplicate_attribute)
2078 ResAttrs.ResourceClass = RC;
2079 HasResourceClass =
true;
2082 case attr::HLSLResourceDimension: {
2083 llvm::dxil::ResourceDimension RD =
2085 if (HasResourceDimension) {
2087 ? diag::warn_duplicate_attribute_exact
2088 : diag::warn_duplicate_attribute)
2092 ResAttrs.ResourceDimension = RD;
2093 HasResourceDimension =
true;
2097 if (ResAttrs.IsROV) {
2101 ResAttrs.IsROV =
true;
2103 case attr::HLSLRawBuffer:
2104 if (ResAttrs.RawBuffer) {
2108 ResAttrs.RawBuffer =
true;
2110 case attr::HLSLIsCounter:
2111 if (ResAttrs.IsCounter) {
2115 ResAttrs.IsCounter =
true;
2117 case attr::HLSLContainedType: {
2120 if (!ContainedTy.
isNull()) {
2122 ? diag::warn_duplicate_attribute_exact
2123 : diag::warn_duplicate_attribute)
2132 llvm_unreachable(
"unhandled resource attribute type");
2136 if (!HasResourceClass) {
2137 S.
Diag(AttrList.back()->getRange().getEnd(),
2138 diag::err_hlsl_missing_resource_class);
2143 Wrapped, ContainedTy, ResAttrs);
2145 if (LocInfo && ContainedTyInfo) {
2158 if (!T->isHLSLResourceType()) {
2159 Diag(AL.
getLoc(), diag::err_hlsl_attribute_needs_intangible_type)
2174 AttributeCommonInfo::AS_CXX11, 0, false ,
2179 case ParsedAttr::AT_HLSLResourceClass: {
2181 Diag(AL.
getLoc(), diag::err_attribute_argument_type)
2192 if (!HLSLResourceClassAttr::ConvertStrToResourceClass(Identifier, RC)) {
2193 Diag(ArgLoc, diag::warn_attribute_type_not_supported)
2194 <<
"ResourceClass" << Identifier;
2197 A = HLSLResourceClassAttr::Create(
getASTContext(), RC, ACI);
2201 case ParsedAttr::AT_HLSLResourceDimension: {
2202 StringRef Identifier;
2204 if (!
SemaRef.checkStringLiteralArgumentAttr(AL, 0, Identifier, &ArgLoc))
2208 llvm::dxil::ResourceDimension RD;
2209 if (!HLSLResourceDimensionAttr::ConvertStrToResourceDimension(Identifier,
2211 Diag(ArgLoc, diag::warn_attribute_type_not_supported)
2212 <<
"ResourceDimension" << Identifier;
2215 A = HLSLResourceDimensionAttr::Create(
getASTContext(), RD, ACI);
2219 case ParsedAttr::AT_HLSLROV:
2223 case ParsedAttr::AT_HLSLRawBuffer:
2227 case ParsedAttr::AT_HLSLIsCounter:
2231 case ParsedAttr::AT_HLSLContainedType: {
2233 Diag(AL.
getLoc(), diag::err_attribute_wrong_number_arguments) << AL << 1;
2239 assert(TSI &&
"no type source info for attribute argument");
2241 diag::err_incomplete_type))
2243 A = HLSLContainedTypeAttr::Create(
getASTContext(), TSI, ACI);
2248 llvm_unreachable(
"unhandled HLSL attribute");
2251 HLSLResourcesTypeAttrs.emplace_back(A);
2257 if (!HLSLResourcesTypeAttrs.size())
2263 HLSLResourcesTypeAttrs, QT, &LocInfo)) {
2264 const HLSLAttributedResourceType *RT =
2271 LocsForHLSLAttributedResources.insert(std::pair(RT, LocInfo));
2273 HLSLResourcesTypeAttrs.clear();
2281 auto I = LocsForHLSLAttributedResources.find(RT);
2282 if (I != LocsForHLSLAttributedResources.end()) {
2283 LocInfo = I->second;
2284 LocsForHLSLAttributedResources.erase(I);
2293void SemaHLSL::collectResourceBindingsOnUserRecordDecl(
const VarDecl *VD,
2294 const RecordType *RT) {
2295 const RecordDecl *RD = RT->getDecl()->getDefinitionOrSelf();
2302 "incomplete arrays inside user defined types are not supported");
2311 if (
const HLSLAttributedResourceType *AttrResType =
2312 HLSLAttributedResourceType::findHandleTypeOnResource(Ty)) {
2317 Bindings.addDeclBindingInfo(VD, RC);
2318 }
else if (
const RecordType *RT = dyn_cast<RecordType>(Ty)) {
2324 collectResourceBindingsOnUserRecordDecl(VD, RT);
2336 bool SpecifiedSpace) {
2337 int RegTypeNum =
static_cast<int>(RegType);
2340 if (D->
hasAttr<HLSLGroupSharedAddressSpaceAttr>()) {
2341 S.
Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
2346 if (
HLSLBufferDecl *CBufferOrTBuffer = dyn_cast<HLSLBufferDecl>(D)) {
2347 ResourceClass RC = CBufferOrTBuffer->isCBuffer() ? ResourceClass::CBuffer
2348 : ResourceClass::SRV;
2358 assert(
isa<VarDecl>(D) &&
"D is expected to be VarDecl or HLSLBufferDecl");
2362 if (
const HLSLAttributedResourceType *AttrResType =
2363 HLSLAttributedResourceType::findHandleTypeOnResource(
2380 if (SpecifiedSpace && !DeclaredInCOrTBuffer)
2381 S.
Diag(ArgLoc, diag::err_hlsl_space_on_global_constant);
2386 if (RegType == RegisterType::CBuffer)
2387 S.
Diag(ArgLoc, diag::warn_hlsl_deprecated_register_type_b);
2388 else if (RegType != RegisterType::C)
2389 S.
Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
2393 if (RegType == RegisterType::C)
2394 S.
Diag(ArgLoc, diag::warn_hlsl_register_type_c_packoffset);
2396 S.
Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
2406 S.
Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
2414 bool RegisterTypesDetected[5] = {
false};
2415 RegisterTypesDetected[
static_cast<int>(regType)] =
true;
2418 if (HLSLResourceBindingAttr *
attr =
2419 dyn_cast<HLSLResourceBindingAttr>(*it)) {
2422 if (RegisterTypesDetected[
static_cast<int>(otherRegType)]) {
2423 int otherRegTypeNum =
static_cast<int>(otherRegType);
2425 diag::err_hlsl_duplicate_register_annotation)
2429 RegisterTypesDetected[
static_cast<int>(otherRegType)] =
true;
2437 bool SpecifiedSpace) {
2442 "expecting VarDecl or HLSLBufferDecl");
2454 const uint64_t &Limit,
2457 uint64_t ArrayCount = 1) {
2462 if (StartSlot > Limit)
2466 if (
const auto *AT = dyn_cast<ArrayType>(T)) {
2469 if (
const auto *CAT = dyn_cast<ConstantArrayType>(AT))
2470 Count = CAT->
getSize().getZExtValue();
2474 ArrayCount * Count);
2478 if (
auto ResTy = dyn_cast<HLSLAttributedResourceType>(T)) {
2481 if (ResTy->getAttrs().ResourceClass != ResClass)
2485 uint64_t EndSlot = StartSlot + ArrayCount - 1;
2486 if (EndSlot > Limit)
2490 StartSlot = EndSlot + 1;
2495 if (
const auto *RT = dyn_cast<RecordType>(T)) {
2498 if (
const auto *CXXRD = dyn_cast<CXXRecordDecl>(RD)) {
2501 ResClass, Ctx, ArrayCount))
2508 ResClass, Ctx, ArrayCount))
2522 const uint64_t Limit = UINT32_MAX;
2523 if (SlotNum > Limit)
2528 if (RegTy == RegisterType::C || RegTy == RegisterType::I)
2531 if (
VarDecl *VD = dyn_cast<VarDecl>(TheDecl)) {
2532 uint64_t BaseSlot = SlotNum;
2540 return (BaseSlot > Limit);
2547 return (SlotNum > Limit);
2550 llvm_unreachable(
"unexpected decl type");
2554 if (
VarDecl *VD = dyn_cast<VarDecl>(TheDecl)) {
2556 if (
const auto *IAT = dyn_cast<IncompleteArrayType>(Ty))
2557 Ty = IAT->getElementType();
2559 diag::err_incomplete_type))
2563 StringRef Slot =
"";
2564 StringRef Space =
"";
2568 Diag(AL.
getLoc(), diag::err_attribute_argument_type)
2578 Diag(AL.
getLoc(), diag::err_attribute_argument_type)
2584 SpaceLoc = Loc->
getLoc();
2587 if (Str.starts_with(
"space")) {
2589 SpaceLoc = Loc->
getLoc();
2598 std::optional<unsigned> SlotNum;
2599 unsigned SpaceNum = 0;
2602 if (!Slot.empty()) {
2604 Diag(SlotLoc, diag::err_hlsl_binding_type_invalid) << Slot.substr(0, 1);
2607 if (RegType == RegisterType::I) {
2608 Diag(SlotLoc, diag::warn_hlsl_deprecated_register_type_i);
2611 const StringRef SlotNumStr = Slot.substr(1);
2616 if (SlotNumStr.getAsInteger(10, N)) {
2617 Diag(SlotLoc, diag::err_hlsl_unsupported_register_number);
2625 Diag(SlotLoc, diag::err_hlsl_register_number_too_large);
2634 if (!Space.starts_with(
"space")) {
2635 Diag(SpaceLoc, diag::err_hlsl_expected_space) << Space;
2638 StringRef SpaceNumStr = Space.substr(5);
2639 if (SpaceNumStr.getAsInteger(10, SpaceNum)) {
2640 Diag(SpaceLoc, diag::err_hlsl_expected_space) << Space;
2645 if (SlotNum.has_value())
2650 HLSLResourceBindingAttr *NewAttr =
2651 HLSLResourceBindingAttr::Create(
getASTContext(), Slot, Space, AL);
2653 NewAttr->setBinding(RegType, SlotNum, SpaceNum);
2708 llvm::DenseMap<const FunctionDecl *, unsigned> ScannedDecls;
2712 llvm::Triple::EnvironmentType CurrentShaderEnvironment;
2713 unsigned CurrentShaderStageBit;
2718 bool ReportOnlyShaderStageIssues;
2721 void SetShaderStageContext(llvm::Triple::EnvironmentType ShaderType) {
2722 static_assert(
sizeof(
unsigned) >= 4);
2723 assert(HLSLShaderAttr::isValidShaderType(ShaderType));
2724 assert((
unsigned)(ShaderType - llvm::Triple::Pixel) < 31 &&
2725 "ShaderType is too big for this bitmap");
2728 unsigned bitmapIndex = ShaderType - llvm::Triple::Pixel;
2729 CurrentShaderEnvironment = ShaderType;
2730 CurrentShaderStageBit = (1 << bitmapIndex);
2733 void SetUnknownShaderStageContext() {
2734 CurrentShaderEnvironment = llvm::Triple::UnknownEnvironment;
2735 CurrentShaderStageBit = (1 << 31);
2738 llvm::Triple::EnvironmentType GetCurrentShaderEnvironment()
const {
2739 return CurrentShaderEnvironment;
2742 bool InUnknownShaderStageContext()
const {
2743 return CurrentShaderEnvironment == llvm::Triple::UnknownEnvironment;
2747 void AddToScannedFunctions(
const FunctionDecl *FD) {
2748 unsigned &ScannedStages = ScannedDecls[FD];
2749 ScannedStages |= CurrentShaderStageBit;
2752 unsigned GetScannedStages(
const FunctionDecl *FD) {
return ScannedDecls[FD]; }
2754 bool WasAlreadyScannedInCurrentStage(
const FunctionDecl *FD) {
2755 return WasAlreadyScannedInCurrentStage(GetScannedStages(FD));
2758 bool WasAlreadyScannedInCurrentStage(
unsigned ScannerStages) {
2759 return ScannerStages & CurrentShaderStageBit;
2762 static bool NeverBeenScanned(
unsigned ScannedStages) {
2763 return ScannedStages == 0;
2767 void HandleFunctionOrMethodRef(FunctionDecl *FD, Expr *RefExpr);
2768 void CheckDeclAvailability(NamedDecl *D,
const AvailabilityAttr *AA,
2770 const AvailabilityAttr *FindAvailabilityAttr(
const Decl *D);
2771 bool HasMatchingEnvironmentOrNone(
const AvailabilityAttr *AA);
2774 DiagnoseHLSLAvailability(Sema &SemaRef)
2776 CurrentShaderEnvironment(llvm::Triple::UnknownEnvironment),
2777 CurrentShaderStageBit(0), ReportOnlyShaderStageIssues(
false) {}
2780 void RunOnTranslationUnit(
const TranslationUnitDecl *TU);
2781 void RunOnFunction(
const FunctionDecl *FD);
2783 bool VisitDeclRefExpr(DeclRefExpr *DRE)
override {
2784 FunctionDecl *FD = llvm::dyn_cast<FunctionDecl>(DRE->
getDecl());
2786 HandleFunctionOrMethodRef(FD, DRE);
2790 bool VisitMemberExpr(MemberExpr *ME)
override {
2791 FunctionDecl *FD = llvm::dyn_cast<FunctionDecl>(ME->
getMemberDecl());
2793 HandleFunctionOrMethodRef(FD, ME);
2798void DiagnoseHLSLAvailability::HandleFunctionOrMethodRef(
FunctionDecl *FD,
2801 "expected DeclRefExpr or MemberExpr");
2805 if (FD->
hasBody(FDWithBody)) {
2806 if (!WasAlreadyScannedInCurrentStage(FDWithBody))
2807 DeclsToScan.push_back(FDWithBody);
2812 const AvailabilityAttr *AA = FindAvailabilityAttr(FD);
2814 CheckDeclAvailability(
2818void DiagnoseHLSLAvailability::RunOnTranslationUnit(
2827 DeclContextsToScan.push_back(TU);
2829 while (!DeclContextsToScan.empty()) {
2830 const DeclContext *DC = DeclContextsToScan.pop_back_val();
2831 for (
auto &D : DC->
decls()) {
2838 if (llvm::dyn_cast<NamespaceDecl>(D) || llvm::dyn_cast<ExportDecl>(D)) {
2839 DeclContextsToScan.push_back(llvm::dyn_cast<DeclContext>(D));
2844 const FunctionDecl *FD = llvm::dyn_cast<FunctionDecl>(D);
2849 if (HLSLShaderAttr *ShaderAttr = FD->
getAttr<HLSLShaderAttr>()) {
2850 SetShaderStageContext(ShaderAttr->getType());
2859 for (
const auto *Redecl : FD->
redecls()) {
2860 if (Redecl->isInExportDeclContext()) {
2867 SetUnknownShaderStageContext();
2875void DiagnoseHLSLAvailability::RunOnFunction(
const FunctionDecl *FD) {
2876 assert(DeclsToScan.empty() &&
"DeclsToScan should be empty");
2877 DeclsToScan.push_back(FD);
2879 while (!DeclsToScan.empty()) {
2887 const unsigned ScannedStages = GetScannedStages(FD);
2888 if (WasAlreadyScannedInCurrentStage(ScannedStages))
2891 ReportOnlyShaderStageIssues = !NeverBeenScanned(ScannedStages);
2893 AddToScannedFunctions(FD);
2898bool DiagnoseHLSLAvailability::HasMatchingEnvironmentOrNone(
2899 const AvailabilityAttr *AA) {
2904 llvm::Triple::EnvironmentType CurrentEnv = GetCurrentShaderEnvironment();
2905 if (CurrentEnv == llvm::Triple::UnknownEnvironment)
2908 llvm::Triple::EnvironmentType AttrEnv =
2909 AvailabilityAttr::getEnvironmentType(IIEnvironment->
getName());
2911 return CurrentEnv == AttrEnv;
2914const AvailabilityAttr *
2915DiagnoseHLSLAvailability::FindAvailabilityAttr(
const Decl *D) {
2916 AvailabilityAttr
const *PartialMatch =
nullptr;
2920 for (
const auto *A : D->
attrs()) {
2921 if (
const auto *Avail = dyn_cast<AvailabilityAttr>(A)) {
2922 StringRef AttrPlatform = Avail->getPlatform()->getName();
2923 StringRef TargetPlatform =
2927 if (AttrPlatform == TargetPlatform) {
2929 if (HasMatchingEnvironmentOrNone(Avail))
2931 PartialMatch = Avail;
2935 return PartialMatch;
2940void DiagnoseHLSLAvailability::CheckDeclAvailability(
NamedDecl *D,
2941 const AvailabilityAttr *AA,
2960 if (ReportOnlyShaderStageIssues)
2966 if (InUnknownShaderStageContext())
2971 bool EnvironmentMatches = HasMatchingEnvironmentOrNone(AA);
2972 VersionTuple Introduced = AA->getIntroduced();
2981 llvm::StringRef PlatformName(
2984 llvm::StringRef CurrentEnvStr =
2985 llvm::Triple::getEnvironmentTypeName(GetCurrentShaderEnvironment());
2987 llvm::StringRef AttrEnvStr =
2988 AA->getEnvironment() ? AA->getEnvironment()->getName() :
"";
2989 bool UseEnvironment = !AttrEnvStr.empty();
2991 if (EnvironmentMatches) {
2992 SemaRef.
Diag(
Range.getBegin(), diag::warn_hlsl_availability)
2993 <<
Range << D << PlatformName << Introduced.getAsString()
2994 << UseEnvironment << CurrentEnvStr;
2996 SemaRef.
Diag(
Range.getBegin(), diag::warn_hlsl_availability_unavailable)
3000 SemaRef.
Diag(D->
getLocation(), diag::note_partial_availability_specified_here)
3001 << D << PlatformName << Introduced.getAsString()
3003 << UseEnvironment << AttrEnvStr << CurrentEnvStr;
3010 if (!DefaultCBufferDecls.empty()) {
3013 DefaultCBufferDecls);
3016 SemaRef.getCurLexicalContext()->addDecl(DefaultCBuffer);
3020 for (
const Decl *VD : DefaultCBufferDecls) {
3021 const HLSLResourceBindingAttr *RBA =
3022 VD->
getAttr<HLSLResourceBindingAttr>();
3023 if (RBA && RBA->hasRegisterSlot() &&
3024 RBA->getRegisterType() == HLSLResourceBindingAttr::RegisterType::C) {
3031 SemaRef.Consumer.HandleTopLevelDecl(DG);
3033 diagnoseAvailabilityViolations(TU);
3043 TI.
getTriple().getEnvironment() != llvm::Triple::EnvironmentType::Library)
3046 DiagnoseHLSLAvailability(
SemaRef).RunOnTranslationUnit(TU);
3053 for (
unsigned I = 1, N = TheCall->
getNumArgs(); I < N; ++I) {
3056 S->
Diag(TheCall->
getBeginLoc(), diag::err_vec_builtin_incompatible_vector)
3081 for (
unsigned I = 0; I < TheCall->
getNumArgs(); ++I) {
3096 if (!BaseType->isFloat32Type())
3097 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3098 << ArgOrdinal << 5 << 0
3110 if (!BaseType->isHalfType() && !BaseType->isFloat32Type())
3111 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3112 << ArgOrdinal << 5 << 0
3118 unsigned ArgIndex) {
3119 auto *Arg = TheCall->
getArg(ArgIndex);
3121 if (Arg->IgnoreCasts()->isModifiableLvalue(S->
Context, &OrigLoc) ==
3124 S->
Diag(OrigLoc, diag::error_hlsl_inout_lvalue) << Arg << 0;
3134 if (VecTy->getElementType()->isDoubleType())
3135 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3136 << ArgOrdinal << 1 << 0 << 1
3146 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3147 << ArgOrdinal << 5 << 1
3156 if (VecTy->getElementType()->isUnsignedIntegerType())
3159 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3160 << ArgOrdinal << 4 << 3 << 0
3169 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3170 << ArgOrdinal << 5 << 3
3176 unsigned ArgOrdinal,
unsigned Width) {
3179 ArgTy = VTy->getElementType();
3181 uint64_t ElementBitCount =
3183 if (ElementBitCount != Width) {
3185 diag::err_integer_incorrect_bit_count)
3186 << Width << ElementBitCount;
3203 unsigned ArgIndex) {
3212 diag::err_typecheck_expect_scalar_or_vector)
3213 << ArgType << Scalar;
3220 QualType Scalar,
unsigned ArgIndex) {
3231 if (
const auto *VTy = ArgType->getAs<
VectorType>()) {
3244 diag::err_typecheck_expect_scalar_or_vector_or_matrix)
3245 << ArgType << Scalar;
3250 unsigned ArgIndex) {
3255 if (!(ArgType->isScalarType() ||
3256 (VTy && VTy->getElementType()->isScalarType()))) {
3258 diag::err_typecheck_expect_any_scalar_or_vector)
3268 unsigned ArgIndex) {
3270 assert(ArgIndex < TheCall->getNumArgs());
3278 diag::err_typecheck_expect_any_scalar_or_vector)
3303 diag::err_typecheck_call_different_arg_types)
3322 Arg1ScalarTy = VTy->getElementType();
3326 Arg2ScalarTy = VTy->getElementType();
3329 S->
Diag(Arg1->
getBeginLoc(), diag::err_hlsl_builtin_scalar_vector_mismatch)
3330 << 1 << TheCall->
getCallee() << Arg1Ty << Arg2Ty;
3340 if (Arg1Length > 0 && Arg0Length != Arg1Length) {
3342 diag::err_typecheck_vector_lengths_not_equal)
3348 if (Arg2Length > 0 && Arg0Length != Arg2Length) {
3350 diag::err_typecheck_vector_lengths_not_equal)
3363 llvm::function_ref<
bool(
const HLSLAttributedResourceType *ResType)> Check =
3367 const HLSLAttributedResourceType *ResTy =
3371 diag::err_typecheck_expect_hlsl_resource)
3375 if (Check && Check(ResTy)) {
3377 diag::err_invalid_hlsl_resource_type)
3385 QualType BaseType,
unsigned ExpectedCount,
3387 unsigned PassedCount = 1;
3389 PassedCount = VecTy->getNumElements();
3391 if (PassedCount != ExpectedCount) {
3394 S->
Diag(Loc, diag::err_typecheck_convert_incompatible)
3406 [](
const HLSLAttributedResourceType *ResType) {
3407 return ResType->getAttrs().ResourceDimension ==
3408 llvm::dxil::ResourceDimension::Unknown;
3414 [](
const HLSLAttributedResourceType *ResType) {
3415 return ResType->getAttrs().ResourceClass !=
3416 llvm::hlsl::ResourceClass::Sampler;
3424 unsigned ExpectedDim =
3441 unsigned NextIdx = 3;
3447 diag::err_typecheck_convert_incompatible)
3455 Expr *ComponentArg = TheCall->
getArg(NextIdx);
3459 diag::err_typecheck_convert_incompatible)
3466 std::optional<llvm::APSInt> ComponentOpt =
3469 int64_t ComponentVal = ComponentOpt->getSExtValue();
3470 if (ComponentVal != 0) {
3473 assert(ComponentVal >= 0 && ComponentVal <= 3 &&
3474 "The component is not in the expected range.");
3476 diag::err_hlsl_gathercmp_invalid_component)
3486 const HLSLAttributedResourceType *ResourceTy =
3489 unsigned ExpectedDim =
3498 assert(ResourceTy->hasContainedType() &&
3499 "Expecting a contained type for resource with a dimension "
3501 QualType ReturnType = ResourceTy->getContainedType();
3505 S.
Diag(TheCall->
getBeginLoc(), diag::err_hlsl_samplecmp_requires_float);
3511 ReturnType = VecTy->getElementType();
3524 [](
const HLSLAttributedResourceType *ResType) {
3525 return ResType->getAttrs().ResourceDimension ==
3526 llvm::dxil::ResourceDimension::Unknown;
3534 unsigned ExpectedDim =
3543 EltTy = VTy->getElementType();
3558 TheCall->
setType(ResourceTy->getContainedType());
3563 unsigned MinArgs, MaxArgs;
3591 const HLSLAttributedResourceType *ResourceTy =
3593 unsigned ExpectedDim =
3596 unsigned NextIdx = 3;
3605 diag::err_typecheck_convert_incompatible)
3640 diag::err_typecheck_convert_incompatible)
3646 assert(ResourceTy->hasContainedType() &&
3647 "Expecting a contained type for resource with a dimension "
3649 QualType ReturnType = ResourceTy->getContainedType();
3652 S.
Diag(TheCall->
getBeginLoc(), diag::err_hlsl_samplecmp_requires_float);
3665 switch (BuiltinID) {
3666 case Builtin::BI__builtin_hlsl_adduint64: {
3667 if (
SemaRef.checkArgCount(TheCall, 2))
3681 if (NumElementsArg != 2 && NumElementsArg != 4) {
3683 << 1 << 64 << NumElementsArg * 32;
3697 case Builtin::BI__builtin_hlsl_resource_getpointer: {
3698 if (
SemaRef.checkArgCount(TheCall, 2) ||
3701 SemaRef.getASTContext().UnsignedIntTy))
3706 QualType ContainedTy = ResourceTy->getContainedType();
3709 ReturnType =
SemaRef.Context.getPointerType(ReturnType);
3715 case Builtin::BI__builtin_hlsl_resource_getpointer_typed: {
3716 if (
SemaRef.checkArgCount(TheCall, 3) ||
3719 SemaRef.getASTContext().UnsignedIntTy))
3724 "expected pointer type for second argument");
3731 diag::err_invalid_use_of_array_type);
3735 ReturnType =
SemaRef.Context.getPointerType(ReturnType);
3740 case Builtin::BI__builtin_hlsl_resource_load_with_status: {
3741 if (
SemaRef.checkArgCount(TheCall, 3) ||
3744 SemaRef.getASTContext().UnsignedIntTy) ||
3746 SemaRef.getASTContext().UnsignedIntTy) ||
3752 QualType ReturnType = ResourceTy->getContainedType();
3757 case Builtin::BI__builtin_hlsl_resource_load_with_status_typed: {
3758 if (
SemaRef.checkArgCount(TheCall, 4) ||
3761 SemaRef.getASTContext().UnsignedIntTy) ||
3763 SemaRef.getASTContext().UnsignedIntTy) ||
3769 "expected pointer type for second argument");
3776 diag::err_invalid_use_of_array_type);
3782 case Builtin::BI__builtin_hlsl_resource_load_level:
3784 case Builtin::BI__builtin_hlsl_resource_sample:
3786 case Builtin::BI__builtin_hlsl_resource_sample_bias:
3788 case Builtin::BI__builtin_hlsl_resource_sample_grad:
3790 case Builtin::BI__builtin_hlsl_resource_sample_level:
3792 case Builtin::BI__builtin_hlsl_resource_sample_cmp:
3794 case Builtin::BI__builtin_hlsl_resource_sample_cmp_level_zero:
3796 case Builtin::BI__builtin_hlsl_resource_gather:
3798 case Builtin::BI__builtin_hlsl_resource_gather_cmp:
3800 case Builtin::BI__builtin_hlsl_resource_uninitializedhandle: {
3801 assert(TheCall->
getNumArgs() == 1 &&
"expected 1 arg");
3807 case Builtin::BI__builtin_hlsl_resource_handlefrombinding: {
3808 assert(TheCall->
getNumArgs() == 6 &&
"expected 6 args");
3814 case Builtin::BI__builtin_hlsl_resource_handlefromimplicitbinding: {
3815 assert(TheCall->
getNumArgs() == 6 &&
"expected 6 args");
3821 case Builtin::BI__builtin_hlsl_resource_counterhandlefromimplicitbinding: {
3822 assert(TheCall->
getNumArgs() == 3 &&
"expected 3 args");
3825 auto *MainResType = MainHandleTy->
getAs<HLSLAttributedResourceType>();
3826 auto MainAttrs = MainResType->getAttrs();
3827 assert(!MainAttrs.IsCounter &&
"cannot create a counter from a counter");
3828 MainAttrs.IsCounter =
true;
3830 MainResType->getWrappedType(), MainResType->getContainedType(),
3834 TheCall->
setType(CounterHandleTy);
3837 case Builtin::BI__builtin_hlsl_and:
3838 case Builtin::BI__builtin_hlsl_or: {
3839 if (
SemaRef.checkArgCount(TheCall, 2))
3853 case Builtin::BI__builtin_hlsl_all:
3854 case Builtin::BI__builtin_hlsl_any: {
3855 if (
SemaRef.checkArgCount(TheCall, 1))
3861 case Builtin::BI__builtin_hlsl_asdouble: {
3862 if (
SemaRef.checkArgCount(TheCall, 2))
3866 SemaRef.Context.UnsignedIntTy,
3871 SemaRef.Context.UnsignedIntTy,
3880 case Builtin::BI__builtin_hlsl_elementwise_clamp: {
3881 if (
SemaRef.BuiltinElementwiseTernaryMath(
3887 case Builtin::BI__builtin_hlsl_dot: {
3889 if (
SemaRef.BuiltinVectorToScalarMath(TheCall))
3895 case Builtin::BI__builtin_hlsl_elementwise_firstbithigh:
3896 case Builtin::BI__builtin_hlsl_elementwise_firstbitlow: {
3897 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
3907 EltTy = VecTy->getElementType();
3908 ResTy =
SemaRef.Context.getExtVectorType(ResTy, VecTy->getNumElements());
3921 case Builtin::BI__builtin_hlsl_select: {
3922 if (
SemaRef.checkArgCount(TheCall, 3))
3930 if (VTy && VTy->getElementType()->isBooleanType() &&
3935 case Builtin::BI__builtin_hlsl_elementwise_saturate:
3936 case Builtin::BI__builtin_hlsl_elementwise_rcp: {
3937 if (
SemaRef.checkArgCount(TheCall, 1))
3943 diag::err_builtin_invalid_arg_type)
3946 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
3950 case Builtin::BI__builtin_hlsl_elementwise_degrees:
3951 case Builtin::BI__builtin_hlsl_elementwise_radians:
3952 case Builtin::BI__builtin_hlsl_elementwise_rsqrt:
3953 case Builtin::BI__builtin_hlsl_elementwise_frac:
3954 case Builtin::BI__builtin_hlsl_elementwise_ddx_coarse:
3955 case Builtin::BI__builtin_hlsl_elementwise_ddy_coarse:
3956 case Builtin::BI__builtin_hlsl_elementwise_ddx_fine:
3957 case Builtin::BI__builtin_hlsl_elementwise_ddy_fine: {
3958 if (
SemaRef.checkArgCount(TheCall, 1))
3963 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
3967 case Builtin::BI__builtin_hlsl_elementwise_isinf:
3968 case Builtin::BI__builtin_hlsl_elementwise_isnan: {
3969 if (
SemaRef.checkArgCount(TheCall, 1))
3974 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
3979 case Builtin::BI__builtin_hlsl_lerp: {
3980 if (
SemaRef.checkArgCount(TheCall, 3))
3987 if (
SemaRef.BuiltinElementwiseTernaryMath(TheCall))
3991 case Builtin::BI__builtin_hlsl_mad: {
3992 if (
SemaRef.BuiltinElementwiseTernaryMath(
3998 case Builtin::BI__builtin_hlsl_mul: {
3999 if (
SemaRef.checkArgCount(TheCall, 2))
4008 if (
const auto *VTy = T->getAs<
VectorType>())
4009 return VTy->getElementType();
4011 return MTy->getElementType();
4015 QualType EltTy0 = getElemType(Ty0);
4024 if (IsVec0 && IsMat1) {
4027 }
else if (IsMat0 && IsVec1) {
4031 assert(IsMat0 && IsMat1);
4041 case Builtin::BI__builtin_hlsl_normalize: {
4042 if (
SemaRef.checkArgCount(TheCall, 1))
4053 case Builtin::BI__builtin_hlsl_transpose: {
4054 if (
SemaRef.checkArgCount(TheCall, 1))
4063 << 1 << 3 << 0 << 0 << ArgTy;
4068 MatTy->getElementType(), MatTy->getNumColumns(), MatTy->getNumRows());
4072 case Builtin::BI__builtin_hlsl_elementwise_sign: {
4073 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
4081 case Builtin::BI__builtin_hlsl_step: {
4082 if (
SemaRef.checkArgCount(TheCall, 2))
4094 case Builtin::BI__builtin_hlsl_wave_active_all_equal: {
4095 if (
SemaRef.checkArgCount(TheCall, 1))
4109 unsigned NumElts = VecTy->getNumElements();
4119 case Builtin::BI__builtin_hlsl_wave_active_max:
4120 case Builtin::BI__builtin_hlsl_wave_active_min:
4121 case Builtin::BI__builtin_hlsl_wave_active_sum:
4122 case Builtin::BI__builtin_hlsl_wave_active_product: {
4123 if (
SemaRef.checkArgCount(TheCall, 1))
4136 case Builtin::BI__builtin_hlsl_wave_active_bit_xor:
4137 case Builtin::BI__builtin_hlsl_wave_active_bit_or: {
4138 if (
SemaRef.checkArgCount(TheCall, 1))
4153 (VTy && VTy->getElementType()->isIntegerType()))) {
4155 diag::err_builtin_invalid_arg_type)
4156 << ArgTyExpr <<
SemaRef.Context.UnsignedIntTy << 1 << 0 << 0;
4166 case Builtin::BI__builtin_elementwise_bitreverse: {
4174 case Builtin::BI__builtin_hlsl_wave_prefix_count_bits: {
4175 if (
SemaRef.checkArgCount(TheCall, 1))
4180 if (!(
ArgType->isScalarType())) {
4182 diag::err_typecheck_expect_any_scalar_or_vector)
4187 if (!(
ArgType->isBooleanType())) {
4189 diag::err_typecheck_expect_any_scalar_or_vector)
4196 case Builtin::BI__builtin_hlsl_wave_read_lane_at: {
4197 if (
SemaRef.checkArgCount(TheCall, 2))
4205 diag::err_typecheck_convert_incompatible)
4206 << ArgTyIndex <<
SemaRef.Context.UnsignedIntTy << 1 << 0 << 0;
4219 case Builtin::BI__builtin_hlsl_wave_get_lane_index: {
4220 if (
SemaRef.checkArgCount(TheCall, 0))
4224 case Builtin::BI__builtin_hlsl_wave_prefix_sum:
4225 case Builtin::BI__builtin_hlsl_wave_prefix_product: {
4226 if (
SemaRef.checkArgCount(TheCall, 1))
4239 case Builtin::BI__builtin_hlsl_quad_read_across_x: {
4240 if (
SemaRef.checkArgCount(TheCall, 1))
4252 case Builtin::BI__builtin_hlsl_elementwise_splitdouble: {
4253 if (
SemaRef.checkArgCount(TheCall, 3))
4268 case Builtin::BI__builtin_hlsl_elementwise_clip: {
4269 if (
SemaRef.checkArgCount(TheCall, 1))
4276 case Builtin::BI__builtin_elementwise_acos:
4277 case Builtin::BI__builtin_elementwise_asin:
4278 case Builtin::BI__builtin_elementwise_atan:
4279 case Builtin::BI__builtin_elementwise_atan2:
4280 case Builtin::BI__builtin_elementwise_ceil:
4281 case Builtin::BI__builtin_elementwise_cos:
4282 case Builtin::BI__builtin_elementwise_cosh:
4283 case Builtin::BI__builtin_elementwise_exp:
4284 case Builtin::BI__builtin_elementwise_exp2:
4285 case Builtin::BI__builtin_elementwise_exp10:
4286 case Builtin::BI__builtin_elementwise_floor:
4287 case Builtin::BI__builtin_elementwise_fmod:
4288 case Builtin::BI__builtin_elementwise_log:
4289 case Builtin::BI__builtin_elementwise_log2:
4290 case Builtin::BI__builtin_elementwise_log10:
4291 case Builtin::BI__builtin_elementwise_pow:
4292 case Builtin::BI__builtin_elementwise_roundeven:
4293 case Builtin::BI__builtin_elementwise_sin:
4294 case Builtin::BI__builtin_elementwise_sinh:
4295 case Builtin::BI__builtin_elementwise_sqrt:
4296 case Builtin::BI__builtin_elementwise_tan:
4297 case Builtin::BI__builtin_elementwise_tanh:
4298 case Builtin::BI__builtin_elementwise_trunc: {
4304 case Builtin::BI__builtin_hlsl_buffer_update_counter: {
4305 assert(TheCall->
getNumArgs() == 2 &&
"expected 2 args");
4306 auto checkResTy = [](
const HLSLAttributedResourceType *ResTy) ->
bool {
4307 return !(ResTy->getAttrs().ResourceClass == ResourceClass::UAV &&
4308 ResTy->getAttrs().RawBuffer && ResTy->hasContainedType());
4313 std::optional<llvm::APSInt> Offset =
4315 if (!Offset.has_value() ||
std::abs(Offset->getExtValue()) != 1) {
4317 diag::err_hlsl_expect_arg_const_int_one_or_neg_one)
4323 case Builtin::BI__builtin_hlsl_elementwise_f16tof32: {
4324 if (
SemaRef.checkArgCount(TheCall, 1))
4335 ArgTy = VTy->getElementType();
4338 diag::err_builtin_invalid_arg_type)
4347 case Builtin::BI__builtin_hlsl_elementwise_f32tof16: {
4348 if (
SemaRef.checkArgCount(TheCall, 1))
4363 WorkList.push_back(BaseTy);
4364 while (!WorkList.empty()) {
4365 QualType T = WorkList.pop_back_val();
4366 T = T.getCanonicalType().getUnqualifiedType();
4367 if (
const auto *AT = dyn_cast<ConstantArrayType>(T)) {
4375 for (uint64_t Ct = 0; Ct < AT->
getZExtSize(); ++Ct)
4376 llvm::append_range(List, ElementFields);
4381 if (
const auto *VT = dyn_cast<VectorType>(T)) {
4382 List.insert(List.end(), VT->getNumElements(), VT->getElementType());
4385 if (
const auto *MT = dyn_cast<ConstantMatrixType>(T)) {
4386 List.insert(List.end(), MT->getNumElementsFlattened(),
4387 MT->getElementType());
4390 if (
const auto *RD = T->getAsCXXRecordDecl()) {
4391 if (RD->isStandardLayout())
4392 RD = RD->getStandardLayoutBaseWithFields();
4396 if (RD->
isUnion() || !RD->isAggregate()) {
4402 for (
const auto *FD : RD->
fields())
4403 if (!FD->isUnnamedBitField())
4404 FieldTypes.push_back(FD->
getType());
4406 std::reverse(FieldTypes.begin(), FieldTypes.end());
4407 llvm::append_range(WorkList, FieldTypes);
4411 if (!RD->isStandardLayout()) {
4413 for (
const auto &
Base : RD->bases())
4414 FieldTypes.push_back(
Base.getType());
4415 std::reverse(FieldTypes.begin(), FieldTypes.end());
4416 llvm::append_range(WorkList, FieldTypes);
4438 if (
SemaRef.Context.getTypeSize(QT) / 8 > 16)
4444 int ArraySize = VT->getNumElements();
4449 QualType ElTy = VT->getElementType();
4453 if (
SemaRef.Context.getTypeSize(QT) / 8 > 16)
4469 if (
SemaRef.getASTContext().hasSameType(T1, T2))
4478 return llvm::equal(T1Types, T2Types,
4480 return SemaRef.IsLayoutCompatible(LHS, RHS);
4489 bool HadError =
false;
4491 for (
unsigned i = 0, e =
New->getNumParams(); i != e; ++i) {
4499 const auto *NDAttr = NewParam->
getAttr<HLSLParamModifierAttr>();
4500 unsigned NSpellingIdx = (NDAttr ? NDAttr->getSpellingListIndex() : 0);
4501 const auto *ODAttr = OldParam->
getAttr<HLSLParamModifierAttr>();
4502 unsigned OSpellingIdx = (ODAttr ? ODAttr->getSpellingListIndex() : 0);
4504 if (NSpellingIdx != OSpellingIdx) {
4506 diag::err_hlsl_param_qualifier_mismatch)
4507 << NDAttr << NewParam;
4523 if (
SemaRef.getASTContext().hasSameUnqualifiedType(SrcTy, DestTy))
4538 llvm_unreachable(
"HLSL doesn't support pointers.");
4541 llvm_unreachable(
"HLSL doesn't support complex types.");
4543 llvm_unreachable(
"HLSL doesn't support fixed point types.");
4545 llvm_unreachable(
"Should have returned before this");
4555 llvm_unreachable(
"HLSL doesn't support complex types.");
4557 llvm_unreachable(
"HLSL doesn't support fixed point types.");
4562 llvm_unreachable(
"HLSL doesn't support pointers.");
4564 llvm_unreachable(
"Should have returned before this");
4570 llvm_unreachable(
"HLSL doesn't support pointers.");
4573 llvm_unreachable(
"HLSL doesn't support fixed point types.");
4577 llvm_unreachable(
"HLSL doesn't support complex types.");
4580 llvm_unreachable(
"Unhandled scalar cast");
4607 for (
unsigned I = 0, Size = DestTypes.size(); I < Size; ++I) {
4608 if (DestTypes[I]->isUnionType())
4640 if (SrcTypes.size() < DestTypes.size())
4643 unsigned SrcSize = SrcTypes.size();
4644 unsigned DstSize = DestTypes.size();
4646 for (I = 0; I < DstSize && I < SrcSize; I++) {
4647 if (SrcTypes[I]->isUnionType() || DestTypes[I]->isUnionType())
4655 for (; I < SrcSize; I++) {
4656 if (SrcTypes[I]->isUnionType())
4663 assert(Param->hasAttr<HLSLParamModifierAttr>() &&
4664 "We should not get here without a parameter modifier expression");
4665 const auto *
Attr = Param->getAttr<HLSLParamModifierAttr>();
4672 << Arg << (IsInOut ? 1 : 0);
4678 QualType Ty = Param->getType().getNonLValueExprType(Ctx);
4685 << Arg << (IsInOut ? 1 : 0);
4697 SemaRef.PerformCopyInitialization(Entity, Param->getBeginLoc(), ArgOpV);
4703 auto *OpV =
new (Ctx)
4708 Res =
SemaRef.ActOnBinOp(
SemaRef.getCurScope(), Param->getBeginLoc(),
4709 tok::equal, ArgOpV, OpV);
4725 "Pointer and reference types cannot be inout or out parameters");
4726 Ty =
SemaRef.getASTContext().getLValueReferenceType(Ty);
4742 for (
const auto *FD : RD->
fields()) {
4746 assert(RD->getNumBases() <= 1 &&
4747 "HLSL doesn't support multiple inheritance");
4748 return RD->getNumBases()
4753 if (
const auto *AT = dyn_cast<ArrayType>(Ty)) {
4754 if (
const auto *CAT = dyn_cast<ConstantArrayType>(AT))
4766 bool IsVKPushConstant = IsVulkan && VD->
hasAttr<HLSLVkPushConstantAttr>();
4771 !VD->
hasAttr<HLSLVkConstantIdAttr>() && !IsVKPushConstant &&
4777 if (
Decl->getType().hasAddressSpace())
4780 if (
Decl->getType()->isDependentType())
4793 llvm::Triple::Vulkan;
4794 if (IsVulkan &&
Decl->
hasAttr<HLSLVkPushConstantAttr>()) {
4795 if (HasDeclaredAPushConstant)
4801 HasDeclaredAPushConstant =
true;
4830 "expected resource type or array of resources");
4840 const HLSLAttributedResourceType *ResHandleTy =
nullptr;
4841 if (
const auto *AT = dyn_cast<ArrayType>(ResTy.
getTypePtr())) {
4842 const auto *CAT = dyn_cast<ConstantArrayType>(AT);
4843 Range = CAT ? CAT->
getSize().getZExtValue() : 0;
4846 ResHandleTy = HLSLAttributedResourceType::findHandleTypeOnResource(
4852 HLSLResourceBindingAttr::CreateImplicit(S.
getASTContext(),
"",
"0", {});
4864 HLSLAssociatedResourceDeclAttr::CreateImplicit(AST, ResDecl));
4874 EmbeddedResourceNameBuilder &NameBuilder);
4884 assert(RD->
getNumBases() <= 1 &&
"HLSL doesn't support multiple inheritance");
4910 }
else if (
const auto *ArrayTy = dyn_cast<ConstantArrayType>(FDTy)) {
4912 "resource arrays should have been already handled");
4931 if (!SubCAT && !ElementRD)
4934 for (
unsigned I = 0, E = CAT->
getSize().getZExtValue(); I < E; ++I) {
4949void SemaHLSL::handleGlobalStructOrArrayOfWithResources(
VarDecl *VD) {
4950 EmbeddedResourceNameBuilder NameBuilder(VD->
getName());
4954 "Expected non-resource struct or array type");
4961 if (
const auto *CAT = dyn_cast<ConstantArrayType>(VDTy)) {
4970 if (
SemaRef.RequireCompleteType(
4973 diag::err_typecheck_decl_incomplete_type)) {
4987 DefaultCBufferDecls.push_back(VD);
4992 collectResourceBindingsOnVarDecl(VD);
4994 if (VD->
hasAttr<HLSLVkConstantIdAttr>())
5006 processExplicitBindingsOnDecl(VD);
5044 handleGlobalStructOrArrayOfWithResources(VD);
5048 if (VD->
hasAttr<HLSLGroupSharedAddressSpaceAttr>())
5057 "expected resource record type");
5073 const char *CreateMethodName;
5075 CreateMethodName = HasCounter ?
"__createFromBindingWithImplicitCounter"
5076 :
"__createFromBinding";
5078 CreateMethodName = HasCounter
5079 ?
"__createFromImplicitBindingWithImplicitCounter"
5080 :
"__createFromImplicitBinding";
5095 Args.push_back(RegSlot);
5103 Args.push_back(OrderId);
5109 Args.push_back(Space);
5113 Args.push_back(RangeSize);
5117 Args.push_back(Index);
5119 StringRef VarName = VD->
getName();
5127 Args.push_back(NameCast);
5135 Args.push_back(CounterId);
5158 SemaRef.CheckCompleteVariableDeclaration(VD);
5164 "expected array of resource records");
5185 lookupMethod(
SemaRef, ResourceDecl,
5186 HasCounter ?
"__createFromBindingWithImplicitCounter"
5187 :
"__createFromBinding",
5191 CreateMethod = lookupMethod(
5193 HasCounter ?
"__createFromImplicitBindingWithImplicitCounter"
5194 :
"__createFromImplicitBinding",
5226std::optional<const DeclBindingInfo *> SemaHLSL::inferGlobalBinding(
Expr *E) {
5227 if (
auto *Ternary = dyn_cast<ConditionalOperator>(E)) {
5228 auto TrueInfo = inferGlobalBinding(Ternary->getTrueExpr());
5229 auto FalseInfo = inferGlobalBinding(Ternary->getFalseExpr());
5230 if (!TrueInfo || !FalseInfo)
5231 return std::nullopt;
5232 if (*TrueInfo != *FalseInfo)
5233 return std::nullopt;
5237 if (
auto *ASE = dyn_cast<ArraySubscriptExpr>(E))
5246 if (
const auto *AttrResType =
5247 HLSLAttributedResourceType::findHandleTypeOnResource(Ty)) {
5249 return Bindings.getDeclBindingInfo(VD, RC);
5256void SemaHLSL::trackLocalResource(
VarDecl *VD,
Expr *E) {
5257 std::optional<const DeclBindingInfo *> ExprBinding = inferGlobalBinding(E);
5260 diag::warn_hlsl_assigning_local_resource_is_not_unique)
5265 if (*ExprBinding ==
nullptr)
5268 auto PrevBinding = Assigns.find(VD);
5269 if (PrevBinding == Assigns.end()) {
5271 Assigns.insert({VD, *ExprBinding});
5276 if (*ExprBinding != PrevBinding->second) {
5278 diag::warn_hlsl_assigning_local_resource_is_not_unique)
5280 SemaRef.Diag(VD->getLocation(), diag::note_var_declared_here) << VD;
5291 "expected LHS to be a resource record or array of resource records");
5292 if (Opc != BO_Assign)
5297 while (
auto *ASE = dyn_cast<ArraySubscriptExpr>(E))
5305 SemaRef.Diag(Loc, diag::err_hlsl_assign_to_global_resource) << VD;
5310 trackLocalResource(VD, RHSExpr);
5318void SemaHLSL::collectResourceBindingsOnVarDecl(
VarDecl *VD) {
5320 "expected global variable that contains HLSL resource");
5323 if (
const HLSLBufferDecl *CBufferOrTBuffer = dyn_cast<HLSLBufferDecl>(VD)) {
5324 Bindings.addDeclBindingInfo(VD, CBufferOrTBuffer->isCBuffer()
5325 ? ResourceClass::CBuffer
5326 : ResourceClass::SRV);
5339 if (
const HLSLAttributedResourceType *AttrResType =
5340 HLSLAttributedResourceType::findHandleTypeOnResource(Ty)) {
5341 Bindings.addDeclBindingInfo(VD, AttrResType->getAttrs().ResourceClass);
5346 if (
const RecordType *RT = dyn_cast<RecordType>(Ty))
5347 collectResourceBindingsOnUserRecordDecl(VD, RT);
5353void SemaHLSL::processExplicitBindingsOnDecl(
VarDecl *VD) {
5356 bool HasBinding =
false;
5357 for (Attr *A : VD->
attrs()) {
5360 if (
auto PA = VD->
getAttr<HLSLVkPushConstantAttr>())
5361 Diag(PA->getLoc(), diag::err_hlsl_attr_incompatible) << A << PA;
5364 HLSLResourceBindingAttr *RBA = dyn_cast<HLSLResourceBindingAttr>(A);
5365 if (!RBA || !RBA->hasRegisterSlot())
5370 assert(RT != RegisterType::I &&
"invalid or obsolete register type should "
5371 "never have an attribute created");
5373 if (RT == RegisterType::C) {
5374 if (Bindings.hasBindingInfoForDecl(VD))
5376 diag::warn_hlsl_user_defined_type_missing_member)
5377 <<
static_cast<int>(RT);
5385 if (DeclBindingInfo *BI = Bindings.getDeclBindingInfo(VD, RC)) {
5390 diag::warn_hlsl_user_defined_type_missing_member)
5391 <<
static_cast<int>(RT);
5399class InitListTransformer {
5403 QualType *DstIt =
nullptr;
5404 Expr **ArgIt =
nullptr;
5410 bool castInitializer(Expr *E) {
5411 assert(DstIt &&
"This should always be something!");
5412 if (DstIt == DestTypes.end()) {
5414 ArgExprs.push_back(E);
5419 DstIt = DestTypes.begin();
5422 Ctx, *DstIt,
false);
5427 ArgExprs.push_back(
Init);
5432 bool buildInitializerListImpl(Expr *E) {
5434 if (
auto *
Init = dyn_cast<InitListExpr>(E)) {
5435 for (
auto *SubInit :
Init->inits())
5436 if (!buildInitializerListImpl(SubInit))
5445 return castInitializer(E);
5447 if (
auto *VecTy = Ty->
getAs<VectorType>()) {
5452 for (uint64_t I = 0; I <
Size; ++I) {
5454 SizeTy, SourceLocation());
5460 if (!castInitializer(ElExpr.
get()))
5465 if (
auto *MTy = Ty->
getAs<ConstantMatrixType>()) {
5466 unsigned Rows = MTy->getNumRows();
5467 unsigned Cols = MTy->getNumColumns();
5468 QualType ElemTy = MTy->getElementType();
5470 for (
unsigned R = 0;
R < Rows; ++
R) {
5471 for (
unsigned C = 0;
C < Cols; ++
C) {
5484 if (!castInitializer(ElExpr.
get()))
5492 if (
auto *ArrTy = dyn_cast<ConstantArrayType>(Ty.
getTypePtr())) {
5496 for (uint64_t I = 0; I <
Size; ++I) {
5498 SizeTy, SourceLocation());
5503 if (!buildInitializerListImpl(ElExpr.
get()))
5510 llvm::SmallVector<CXXRecordDecl *> RecordDecls;
5511 RecordDecls.push_back(RD);
5512 while (RecordDecls.back()->getNumBases()) {
5513 CXXRecordDecl *D = RecordDecls.back();
5515 "HLSL doesn't support multiple inheritance");
5516 RecordDecls.push_back(
5519 while (!RecordDecls.empty()) {
5520 CXXRecordDecl *RD = RecordDecls.pop_back_val();
5521 for (
auto *FD : RD->
fields()) {
5522 if (FD->isUnnamedBitField())
5530 if (!buildInitializerListImpl(Res.
get()))
5538 Expr *generateInitListsImpl(QualType Ty) {
5539 assert(ArgIt != ArgExprs.end() &&
"Something is off in iteration!");
5543 llvm::SmallVector<Expr *>
Inits;
5549 if (
auto *ATy = Ty->
getAs<VectorType>()) {
5550 ElTy = ATy->getElementType();
5551 Size = ATy->getNumElements();
5552 }
else if (
auto *CMTy = Ty->
getAs<ConstantMatrixType>()) {
5553 ElTy = CMTy->getElementType();
5554 Size = CMTy->getNumElementsFlattened();
5557 ElTy = VTy->getElementType();
5558 Size = VTy->getZExtSize();
5560 for (uint64_t I = 0; I <
Size; ++I)
5561 Inits.push_back(generateInitListsImpl(ElTy));
5564 llvm::SmallVector<CXXRecordDecl *> RecordDecls;
5565 RecordDecls.push_back(RD);
5566 while (RecordDecls.back()->getNumBases()) {
5567 CXXRecordDecl *D = RecordDecls.back();
5569 "HLSL doesn't support multiple inheritance");
5570 RecordDecls.push_back(
5573 while (!RecordDecls.empty()) {
5574 CXXRecordDecl *RD = RecordDecls.pop_back_val();
5575 for (
auto *FD : RD->
fields())
5576 if (!FD->isUnnamedBitField())
5580 auto *NewInit =
new (Ctx) InitListExpr(Ctx,
Inits.front()->getBeginLoc(),
5582 NewInit->setType(Ty);
5587 llvm::SmallVector<QualType, 16> DestTypes;
5588 llvm::SmallVector<Expr *, 16> ArgExprs;
5589 InitListTransformer(Sema &SemaRef,
const InitializedEntity &Entity)
5590 : S(SemaRef), Ctx(SemaRef.getASTContext()),
5591 Wrap(Entity.
getType()->isIncompleteArrayType()) {
5592 InitTy = Entity.
getType().getNonReferenceType();
5602 DstIt = DestTypes.begin();
5605 bool buildInitializerList(Expr *E) {
return buildInitializerListImpl(E); }
5607 Expr *generateInitLists() {
5608 assert(!ArgExprs.empty() &&
5609 "Call buildInitializerList to generate argument expressions.");
5610 ArgIt = ArgExprs.begin();
5612 return generateInitListsImpl(InitTy);
5613 llvm::SmallVector<Expr *>
Inits;
5614 while (ArgIt != ArgExprs.end())
5615 Inits.push_back(generateInitListsImpl(InitTy));
5617 auto *NewInit =
new (Ctx) InitListExpr(Ctx,
Inits.front()->getBeginLoc(),
5619 llvm::APInt ArySize(64,
Inits.size());
5621 ArraySizeModifier::Normal, 0));
5633 if (
const ArrayType *AT = dyn_cast<ArrayType>(Ty)) {
5640 if (
const auto *RT = Ty->
getAs<RecordType>()) {
5644 if (
const auto *CXXRD = dyn_cast<CXXRecordDecl>(RD)) {
5664 if (
Init->getType()->isScalarType())
5667 InitListTransformer ILT(
SemaRef, Entity);
5669 for (
unsigned I = 0; I <
Init->getNumInits(); ++I) {
5677 Init->setInit(I, E);
5679 if (!ILT.buildInitializerList(E))
5682 size_t ExpectedSize = ILT.DestTypes.size();
5683 size_t ActualSize = ILT.ArgExprs.size();
5684 if (ExpectedSize == 0 && ActualSize == 0)
5691 InitTy =
SemaRef.getASTContext().removeAddrSpaceQualType(InitTy);
5693 SemaRef.Diag(
Init->getBeginLoc(), diag::err_hlsl_incorrect_num_initializers)
5694 << (int)(ExpectedSize < ActualSize) << InitTy
5695 << ExpectedSize << ActualSize;
5705 assert(ExpectedSize > 0 &&
5706 "The expected size of an incomplete array type must be at least 1.");
5708 ((ActualSize + ExpectedSize - 1) / ExpectedSize) * ExpectedSize;
5716 InitTy =
SemaRef.getASTContext().removeAddrSpaceQualType(InitTy);
5717 if (ExpectedSize != ActualSize) {
5718 int TooManyOrFew = ActualSize > ExpectedSize ? 1 : 0;
5719 SemaRef.Diag(
Init->getBeginLoc(), diag::err_hlsl_incorrect_num_initializers)
5720 << TooManyOrFew << InitTy << ExpectedSize << ActualSize;
5727 Init->resizeInits(Ctx, NewInit->getNumInits());
5728 for (
unsigned I = 0; I < NewInit->getNumInits(); ++I)
5729 Init->updateInit(Ctx, I, NewInit->getInit(I));
5737 S.
Diag(OpLoc, diag::err_builtin_matrix_invalid_member)
5747 StringRef AccessorName = CompName->
getName();
5748 assert(!AccessorName.empty() &&
"Matrix Accessor must have a name");
5750 unsigned Rows = MT->getNumRows();
5751 unsigned Cols = MT->getNumColumns();
5752 bool IsZeroBasedAccessor =
false;
5753 unsigned ChunkLen = 0;
5754 if (AccessorName.size() < 2)
5756 "length 4 for zero based: \'_mRC\' or "
5757 "length 3 for one-based: \'_RC\' accessor",
5760 if (AccessorName[0] ==
'_') {
5761 if (AccessorName[1] ==
'm') {
5762 IsZeroBasedAccessor =
true;
5769 S, AccessorName,
"zero based: \'_mRC\' or one-based: \'_RC\' accessor",
5772 if (AccessorName.size() % ChunkLen != 0) {
5773 const llvm::StringRef
Expected = IsZeroBasedAccessor
5774 ?
"zero based: '_mRC' accessor"
5775 :
"one-based: '_RC' accessor";
5780 auto isDigit = [](
char c) {
return c >=
'0' &&
c <=
'9'; };
5781 auto isZeroBasedIndex = [](
unsigned i) {
return i <= 3; };
5782 auto isOneBasedIndex = [](
unsigned i) {
return i >= 1 && i <= 4; };
5784 bool HasRepeated =
false;
5786 unsigned NumComponents = 0;
5787 const char *Begin = AccessorName.data();
5789 for (
unsigned I = 0, E = AccessorName.size(); I < E; I += ChunkLen) {
5790 const char *Chunk = Begin + I;
5791 char RowChar = 0, ColChar = 0;
5792 if (IsZeroBasedAccessor) {
5794 if (Chunk[0] !=
'_' || Chunk[1] !=
'm') {
5795 char Bad = (Chunk[0] !=
'_') ? Chunk[0] : Chunk[1];
5797 S, StringRef(&Bad, 1),
"\'_m\' prefix",
5804 if (Chunk[0] !=
'_')
5806 S, StringRef(&Chunk[0], 1),
"\'_\' prefix",
5813 bool IsDigitsError =
false;
5815 unsigned BadPos = IsZeroBasedAccessor ? 2 : 1;
5819 IsDigitsError =
true;
5823 unsigned BadPos = IsZeroBasedAccessor ? 3 : 2;
5827 IsDigitsError =
true;
5832 unsigned Row = RowChar -
'0';
5833 unsigned Col = ColChar -
'0';
5835 bool HasIndexingError =
false;
5836 if (IsZeroBasedAccessor) {
5838 if (!isZeroBasedIndex(Row)) {
5839 S.
Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
5841 HasIndexingError =
true;
5843 if (!isZeroBasedIndex(Col)) {
5844 S.
Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
5846 HasIndexingError =
true;
5850 if (!isOneBasedIndex(Row)) {
5851 S.
Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
5853 HasIndexingError =
true;
5855 if (!isOneBasedIndex(Col)) {
5856 S.
Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
5858 HasIndexingError =
true;
5865 if (HasIndexingError)
5871 bool HasBoundsError =
false;
5873 Diag(OpLoc, diag::err_hlsl_matrix_index_out_of_bounds)
5875 HasBoundsError =
true;
5878 Diag(OpLoc, diag::err_hlsl_matrix_index_out_of_bounds)
5880 HasBoundsError =
true;
5885 unsigned FlatIndex = Row * Cols + Col;
5886 if (Seen[FlatIndex])
5888 Seen[FlatIndex] =
true;
5891 if (NumComponents == 0 || NumComponents > 4) {
5892 S.
Diag(OpLoc, diag::err_hlsl_matrix_swizzle_invalid_length)
5897 QualType ElemTy = MT->getElementType();
5898 if (NumComponents == 1)
5904 for (Sema::ExtVectorDeclsType::iterator
5908 if ((*I)->getUnderlyingType() == VT)
5919 trackLocalResource(VDecl,
Init);
5921 const HLSLVkConstantIdAttr *ConstIdAttr =
5922 VDecl->
getAttr<HLSLVkConstantIdAttr>();
5929 if (!
Init->isCXX11ConstantExpr(Context, &InitValue)) {
5939 int ConstantID = ConstIdAttr->getId();
5940 llvm::APInt IDVal(Context.getIntWidth(Context.IntTy), ConstantID);
5942 ConstIdAttr->getLocation());
5946 if (
C->getType()->getCanonicalTypeUnqualified() !=
5950 Context.getTrivialTypeSourceInfo(
5951 Init->getType(),
Init->getExprLoc()),
5970 if (!Params || Params->
size() != 1)
5983 if (
auto *TTP = dyn_cast<TemplateTypeParmDecl>(P)) {
5984 if (TTP->hasDefaultArgument()) {
5985 TemplateArgs.
addArgument(TTP->getDefaultArgument());
5988 }
else if (
auto *NTTP = dyn_cast<NonTypeTemplateParmDecl>(P)) {
5989 if (NTTP->hasDefaultArgument()) {
5990 TemplateArgs.
addArgument(NTTP->getDefaultArgument());
5993 }
else if (
auto *TTPD = dyn_cast<TemplateTemplateParmDecl>(P)) {
5994 if (TTPD->hasDefaultArgument()) {
5995 TemplateArgs.
addArgument(TTPD->getDefaultArgument());
6002 return SemaRef.CheckTemplateIdType(
6004 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 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 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 handleStructWithResources(Sema &S, VarDecl *ParentVD, const CXXRecordDecl *RD, EmbeddedResourceNameBuilder &NameBuilder)
static void createGlobalResourceDeclForStruct(Sema &S, VarDecl *ParentVD, SourceLocation Loc, IdentifierInfo *Id, QualType ResTy)
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 bool CheckNotBoolScalarOrVector(Sema *S, CallExpr *TheCall, unsigned ArgIndex)
static CXXRecordDecl * findRecordDeclInContext(IdentifierInfo *II, DeclContext *DC)
static bool CheckWavePrefix(Sema *S, CallExpr *TheCall)
static bool CheckExpectedBitWidth(Sema *S, CallExpr *TheCall, unsigned ArgOrdinal, unsigned Width)
static bool hasCounterHandle(const CXXRecordDecl *RD)
static bool CheckVectorSelect(Sema *S, CallExpr *TheCall)
static QualType handleFloatVectorBinOpConversion(Sema &SemaRef, ExprResult &LHS, ExprResult &RHS, QualType LHSType, QualType RHSType, QualType LElTy, QualType RElTy, bool IsCompAssign)
static ResourceClass getResourceClass(RegisterType RT)
static CXXRecordDecl * createHostLayoutStruct(Sema &S, CXXRecordDecl *StructDecl)
static bool CheckScalarOrVector(Sema *S, CallExpr *TheCall, QualType Scalar, unsigned ArgIndex)
static bool CheckSamplingBuiltin(Sema &S, CallExpr *TheCall, SampleKind Kind)
static bool CheckScalarOrVectorOrMatrix(Sema *S, CallExpr *TheCall, QualType Scalar, unsigned ArgIndex)
static bool CheckFloatRepresentation(Sema *S, SourceLocation Loc, int ArgOrdinal, clang::QualType PassedType)
static void handleArrayOfStructWithResources(Sema &S, VarDecl *ParentVD, const ConstantArrayType *CAT, EmbeddedResourceNameBuilder &NameBuilder)
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)
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.
QualType getSizeType() const
Return the unique type for "size_t" (C99 7.17), defined in <stddef.h>.
QualType getExtVectorType(QualType VectorType, unsigned NumElts) const
Return the unique reference to an extended vector type of the specified element type and size.
const TargetInfo & getTargetInfo() const
QualType getHLSLAttributedResourceType(QualType Wrapped, QualType Contained, const HLSLAttributedResourceType::Attributes &Attrs)
QualType getAddrSpaceQualType(QualType T, LangAS AddressSpace) const
Return the uniqued reference to the type for an address space qualified type with the specified type ...
CanQualType getCanonicalTagType(const TagDecl *TD) const
static bool hasSameUnqualifiedType(QualType T1, QualType T2)
Determine whether the given types are equivalent after cvr-qualifiers have been removed.
QualType getConstantMatrixType(QualType ElementType, unsigned NumRows, unsigned NumColumns) const
Return the unique reference to the matrix type of the specified element type and size.
unsigned getTypeAlign(QualType T) const
Return the ABI-specified alignment of a (complete) type T, in bits.
Represents an array type, per C99 6.7.5.2 - Array Declarators.
QualType getElementType() const
Attr - This represents one attribute.
attr::Kind getKind() const
SourceLocation getLocation() const
SourceLocation getScopeLoc() const
SourceRange getRange() const
const IdentifierInfo * getScopeName() const
SourceLocation getLoc() const
const IdentifierInfo * getAttrName() const
Represents a base class of a C++ class.
QualType getType() const
Retrieves the type of the base class.
Represents a static or instance method of a struct/union/class.
Represents a C++ struct/union/class.
bool isHLSLIntangible() const
Returns true if the class contains HLSL intangible type, either as a field or in base class.
static CXXRecordDecl * Create(const ASTContext &C, TagKind TK, DeclContext *DC, SourceLocation StartLoc, SourceLocation IdLoc, IdentifierInfo *Id, CXXRecordDecl *PrevDecl=nullptr)
void setBases(CXXBaseSpecifier const *const *Bases, unsigned NumBases)
Sets the base classes of this struct or class.
base_class_iterator bases_end()
void completeDefinition() override
Indicates that the definition of this class is now complete.
unsigned getNumBases() const
Retrieves the number of base classes of this class.
base_class_iterator bases_begin()
bool isEmpty() const
Determine whether this is an empty class in the sense of (C++11 [meta.unary.prop]).
CallExpr - Represents a function call (C99 6.5.2.2, C++ [expr.call]).
Expr * getArg(unsigned Arg)
getArg - Return the specified argument.
SourceLocation getBeginLoc() const
static CallExpr * Create(const ASTContext &Ctx, Expr *Fn, ArrayRef< Expr * > Args, QualType Ty, ExprValueKind VK, SourceLocation RParenLoc, FPOptionsOverride FPFeatures, unsigned MinNumArgs=0, ADLCallKind UsesADL=NotADL)
Create a call expression.
FunctionDecl * getDirectCallee()
If the callee is a FunctionDecl, return it. Otherwise return null.
unsigned getNumArgs() const
getNumArgs - Return the number of actual arguments to this call.
QualType withConst() const
Retrieves a version of this type with const applied.
const T * getTypePtr() const
Retrieve the underlying type pointer, which refers to a canonical type.
QuantityType getQuantity() const
getQuantity - Get the raw integer representation of this quantity.
Represents the canonical version of C arrays with a specified constant size.
bool isZeroSize() const
Return true if the size is zero.
llvm::APInt getSize() const
Return the constant array size as an APInt.
uint64_t getZExtSize() const
Return the size zero-extended as a uint64_t.
Represents a concrete matrix type with constant number of rows and columns.
unsigned getNumColumns() const
Returns the number of columns in the matrix.
static DeclAccessPair make(NamedDecl *D, AccessSpecifier AS)
DeclContext - This is used only as base class of specific decl types that can act as declaration cont...
lookup_result lookup(DeclarationName Name) const
lookup - Find the declarations (if any) with the given Name in this context.
bool isTranslationUnit() const
void addDecl(Decl *D)
Add the declaration D into this context.
decl_range decls() const
decls_begin/decls_end - Iterate over the declarations stored in this context.
DeclContext * getNonTransparentContext()
A reference to a declared variable, function, enum, etc.
static DeclRefExpr * Create(const ASTContext &Context, NestedNameSpecifierLoc QualifierLoc, SourceLocation TemplateKWLoc, ValueDecl *D, bool RefersToEnclosingVariableOrCapture, SourceLocation NameLoc, QualType T, ExprValueKind VK, NamedDecl *FoundD=nullptr, const TemplateArgumentListInfo *TemplateArgs=nullptr, NonOdrUseReason NOUR=NOUR_None)
Decl - This represents one declaration (or definition), e.g.
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.
ExprValueKind getValueKind() const
getValueKind - The value kind that this expression produces.
Expr * IgnoreParenImpCasts() LLVM_READONLY
Skip past any parentheses and implicit casts which might surround this expression until reaching a fi...
Expr * IgnoreParens() LLVM_READONLY
Skip past any parentheses which might surround this expression until reaching a fixed point.
std::optional< llvm::APSInt > getIntegerConstantExpr(const ASTContext &Ctx) const
isIntegerConstantExpr - Return the value if this expression is a valid integer constant expression.
bool isLValue() const
isLValue - True if this expression is an "l-value" according to the rules of the current language.
ExprObjectKind getObjectKind() const
getObjectKind - The object kind that this expression produces.
bool HasSideEffects(const ASTContext &Ctx, bool IncludePossibleEffects=true) const
HasSideEffects - This routine returns true for all those expressions which have any effect other than...
void setValueKind(ExprValueKind Cat)
setValueKind - Set the value kind produced by this expression.
SourceLocation getExprLoc() const LLVM_READONLY
getExprLoc - Return the preferred location for the arrow when diagnosing a problem with a generic exp...
ExtVectorType - Extended vector type.
Represents difference between two FPOptions values.
Represents a member of a struct/union/class.
static FieldDecl * Create(const ASTContext &C, DeclContext *DC, SourceLocation StartLoc, SourceLocation IdLoc, const IdentifierInfo *Id, QualType T, TypeSourceInfo *TInfo, Expr *BW, bool Mutable, InClassInitStyle InitStyle)
static FixItHint CreateReplacement(CharSourceRange RemoveRange, StringRef Code)
Create a code modification hint that replaces the given source range with the given code string.
Represents a function declaration or definition.
const ParmVarDecl * getParamDecl(unsigned i) const
Stmt * getBody(const FunctionDecl *&Definition) const
Retrieve the body (definition) of the function.
bool isThisDeclarationADefinition() const
Returns whether this specific declaration of the function is also a definition that does not contain ...
QualType getReturnType() const
ArrayRef< ParmVarDecl * > parameters() const
bool isTemplateInstantiation() const
Determines if the given function was instantiated from a function template.
redecl_range redecls() const
Returns an iterator range for all the redeclarations of the same decl.
unsigned getNumParams() const
Return the number of parameters this function must have based on its FunctionType.
DeclarationNameInfo getNameInfo() const
bool hasBody(const FunctionDecl *&Definition) const
Returns true if the function has a body.
bool isDefined(const FunctionDecl *&Definition, bool CheckForPendingFriendDefinition=false) const
Returns true if the function has a definition that does not need to be instantiated.
HLSLBufferDecl - Represent a cbuffer or tbuffer declaration.
static HLSLBufferDecl * Create(ASTContext &C, DeclContext *LexicalParent, bool CBuffer, SourceLocation KwLoc, IdentifierInfo *ID, SourceLocation IDLoc, SourceLocation LBrace)
void addLayoutStruct(CXXRecordDecl *LS)
void setHasValidPackoffset(bool PO)
static HLSLBufferDecl * CreateDefaultCBuffer(ASTContext &C, DeclContext *LexicalParent, ArrayRef< Decl * > DefaultCBufferDecls)
buffer_decl_range buffer_decls() const
static HLSLOutArgExpr * Create(const ASTContext &C, QualType Ty, OpaqueValueExpr *Base, OpaqueValueExpr *OpV, Expr *WB, bool IsInOut)
static HLSLRootSignatureDecl * Create(ASTContext &C, DeclContext *DC, SourceLocation Loc, IdentifierInfo *ID, llvm::dxbc::RootSignatureVersion Version, ArrayRef< llvm::hlsl::rootsig::RootElement > RootElements)
One of these records is kept for each identifier that is lexed.
StringRef getName() const
Return the actual identifier string.
A simple pair of identifier info and location.
SourceLocation getLoc() const
IdentifierInfo * getIdentifierInfo() const
IdentifierInfo & get(StringRef Name)
Return the identifier token info for the specified named identifier.
ImplicitCastExpr - Allows us to explicitly represent implicit type conversions, which have no direct ...
static ImplicitCastExpr * Create(const ASTContext &Context, QualType T, CastKind Kind, Expr *Operand, const CXXCastPath *BasePath, ExprValueKind Cat, FPOptionsOverride FPO)
Describes an C or C++ initializer list.
Describes an entity that is being initialized.
QualType getType() const
Retrieve type being initialized.
static InitializedEntity InitializeParameter(ASTContext &Context, ParmVarDecl *Parm)
Create the initialization entity for a parameter.
static IntegerLiteral * Create(const ASTContext &C, const llvm::APInt &V, QualType type, SourceLocation l)
Returns a new integer literal with value 'V' and type 'type'.
iterator begin(Source *source, bool LocalOnly=false)
Represents the results of name lookup.
Represents a prvalue temporary that is written into memory so that a reference can bind to it.
ValueDecl * getMemberDecl() const
Retrieve the member declaration to which this expression refers.
This represents a decl that may have a name.
IdentifierInfo * getIdentifier() const
Get the identifier that names this declaration, if there is one.
StringRef getName() const
Get the name of identifier for this declaration as a StringRef.
DeclarationName getDeclName() const
Get the actual, stored name of the declaration, which may be a special name.
A C++ nested-name-specifier augmented with source location information.
OpaqueValueExpr - An expression referring to an opaque object of a fixed type and value class.
Represents a parameter to a function.
ParsedAttr - Represents a syntactic attribute.
unsigned getSemanticSpelling() const
If the parsed attribute has a semantic equivalent, and it would have a semantic Spelling enumeration ...
unsigned getMinArgs() const
bool checkExactlyNumArgs(class Sema &S, unsigned Num) const
Check if the attribute has exactly as many args as Num.
IdentifierLoc * getArgAsIdent(unsigned Arg) const
bool hasParsedType() const
const ParsedType & getTypeArg() const
unsigned getNumArgs() const
getNumArgs - Return the number of actual arguments to this attribute.
bool isArgIdent(unsigned Arg) const
Expr * getArgAsExpr(unsigned Arg) const
AttributeCommonInfo::Kind getKind() const
A (possibly-)qualified type.
void addRestrict()
Add the restrict qualifier to this QualType.
QualType getNonLValueExprType(const ASTContext &Context) const
Determine the type of a (typically non-lvalue) expression with the specified result type.
QualType getDesugaredType(const ASTContext &Context) const
Return the specified type with any "sugar" removed from the type.
bool isNull() const
Return true if this QualType doesn't point to a type yet.
const Type * getTypePtr() const
Retrieves a pointer to the underlying (unqualified) type.
LangAS getAddressSpace() const
Return the address space of this type.
QualType getNonReferenceType() const
If Type is a reference type (e.g., const int&), returns the type that the reference refers to ("const...
QualType getCanonicalType() const
QualType getUnqualifiedType() const
Retrieve the unqualified variant of the given type, removing as little sugar as possible.
bool hasAddressSpace() const
Check if this type has any address space qualifier.
Represents a struct/union/class.
field_iterator field_end() const
field_range fields() const
field_iterator field_begin() const
bool hasBindingInfoForDecl(const VarDecl *VD) const
DeclBindingInfo * getDeclBindingInfo(const VarDecl *VD, ResourceClass ResClass)
DeclBindingInfo * addDeclBindingInfo(const VarDecl *VD, ResourceClass ResClass)
Scope - A scope is a transient data structure that is used while parsing the program.
ASTContext & getASTContext() const
SemaDiagnosticBuilder Diag(SourceLocation Loc, unsigned DiagID)
Emit a diagnostic.
ExprResult ActOnOutParamExpr(ParmVarDecl *Param, Expr *Arg)
HLSLRootSignatureDecl * lookupRootSignatureOverrideDecl(DeclContext *DC) const
bool CanPerformElementwiseCast(Expr *Src, QualType DestType)
void handleWaveSizeAttr(Decl *D, const ParsedAttr &AL)
void handleVkLocationAttr(Decl *D, const ParsedAttr &AL)
HLSLAttributedResourceLocInfo TakeLocForHLSLAttribute(const HLSLAttributedResourceType *RT)
void handleSemanticAttr(Decl *D, const ParsedAttr &AL)
bool CanPerformScalarCast(QualType SrcTy, QualType DestTy)
QualType ProcessResourceTypeAttributes(QualType Wrapped)
void handleShaderAttr(Decl *D, const ParsedAttr &AL)
uint32_t getNextImplicitBindingOrderID()
void CheckEntryPoint(FunctionDecl *FD)
void emitLogicalOperatorFixIt(Expr *LHS, Expr *RHS, BinaryOperatorKind Opc)
T * createSemanticAttr(const AttributeCommonInfo &ACI, std::optional< unsigned > Location)
bool initGlobalResourceDecl(VarDecl *VD)
void ActOnEndOfTranslationUnit(TranslationUnitDecl *TU)
bool initGlobalResourceArrayDecl(VarDecl *VD)
HLSLVkConstantIdAttr * mergeVkConstantIdAttr(Decl *D, const AttributeCommonInfo &AL, int Id)
HLSLNumThreadsAttr * mergeNumThreadsAttr(Decl *D, const AttributeCommonInfo &AL, int X, int Y, int Z)
void deduceAddressSpace(VarDecl *Decl)
std::pair< IdentifierInfo *, bool > ActOnStartRootSignatureDecl(StringRef Signature)
Computes the unique Root Signature identifier from the given signature, then lookup if there is a pre...
void handlePackOffsetAttr(Decl *D, const ParsedAttr &AL)
bool diagnosePositionType(QualType T, const ParsedAttr &AL)
bool handleInitialization(VarDecl *VDecl, Expr *&Init)
bool diagnoseInputIDType(QualType T, const ParsedAttr &AL)
void handleParamModifierAttr(Decl *D, const ParsedAttr &AL)
bool CheckResourceBinOp(BinaryOperatorKind Opc, Expr *LHSExpr, Expr *RHSExpr, SourceLocation Loc)
bool CanPerformAggregateSplatCast(Expr *Src, QualType DestType)
bool IsScalarizedLayoutCompatible(QualType T1, QualType T2) const
QualType ActOnTemplateShorthand(TemplateDecl *Template, SourceLocation NameLoc)
void diagnoseSystemSemanticAttr(Decl *D, const ParsedAttr &AL, std::optional< unsigned > Index)
void handleRootSignatureAttr(Decl *D, const ParsedAttr &AL)
bool CheckCompatibleParameterABI(FunctionDecl *New, FunctionDecl *Old)
QualType handleVectorBinOpConversion(ExprResult &LHS, ExprResult &RHS, QualType LHSType, QualType RHSType, bool IsCompAssign)
QualType checkMatrixComponent(Sema &S, QualType baseType, ExprValueKind &VK, SourceLocation OpLoc, const IdentifierInfo *CompName, SourceLocation CompLoc)
void handleResourceBindingAttr(Decl *D, const ParsedAttr &AL)
bool IsTypedResourceElementCompatible(QualType T1)
bool transformInitList(const InitializedEntity &Entity, InitListExpr *Init)
void handleNumThreadsAttr(Decl *D, const ParsedAttr &AL)
bool ActOnUninitializedVarDecl(VarDecl *D)
void handleVkExtBuiltinInputAttr(Decl *D, const ParsedAttr &AL)
void ActOnTopLevelFunction(FunctionDecl *FD)
bool handleResourceTypeAttr(QualType T, const ParsedAttr &AL)
void handleVkPushConstantAttr(Decl *D, const ParsedAttr &AL)
HLSLShaderAttr * mergeShaderAttr(Decl *D, const AttributeCommonInfo &AL, llvm::Triple::EnvironmentType ShaderType)
void ActOnFinishBuffer(Decl *Dcl, SourceLocation RBrace)
void handleVkBindingAttr(Decl *D, const ParsedAttr &AL)
HLSLParamModifierAttr * mergeParamModifierAttr(Decl *D, const AttributeCommonInfo &AL, HLSLParamModifierAttr::Spelling Spelling)
QualType getInoutParameterType(QualType Ty)
void handleVkConstantIdAttr(Decl *D, const ParsedAttr &AL)
Decl * ActOnStartBuffer(Scope *BufferScope, bool CBuffer, SourceLocation KwLoc, IdentifierInfo *Ident, SourceLocation IdentLoc, SourceLocation LBrace)
HLSLWaveSizeAttr * mergeWaveSizeAttr(Decl *D, const AttributeCommonInfo &AL, int Min, int Max, int Preferred, int SpelledArgsCount)
bool handleRootSignatureElements(ArrayRef< hlsl::RootSignatureElement > Elements)
void ActOnFinishRootSignatureDecl(SourceLocation Loc, IdentifierInfo *DeclIdent, ArrayRef< hlsl::RootSignatureElement > Elements)
Creates the Root Signature decl of the parsed Root Signature elements onto the AST and push it onto c...
void ActOnVariableDeclarator(VarDecl *VD)
bool CheckBuiltinFunctionCall(unsigned BuiltinID, CallExpr *TheCall)
Sema - This implements semantic analysis and AST building for C.
@ LookupOrdinaryName
Ordinary name lookup, which finds ordinary names (functions, variables, typedefs, etc....
@ LookupMemberName
Member name lookup, which finds the names of class/struct/union members.
ExtVectorDeclsType ExtVectorDecls
ExtVectorDecls - This is a list all the extended vector types.
ASTContext & getASTContext() const
ExprResult ImpCastExprToType(Expr *E, QualType Type, CastKind CK, ExprValueKind VK=VK_PRValue, const CXXCastPath *BasePath=nullptr, CheckedConversionKind CCK=CheckedConversionKind::Implicit)
ImpCastExprToType - If Expr is not of type 'Type', insert an implicit cast.
const LangOptions & getLangOpts() const
ExprResult BuildFieldReferenceExpr(Expr *BaseExpr, bool IsArrow, SourceLocation OpLoc, const CXXScopeSpec &SS, FieldDecl *Field, DeclAccessPair FoundDecl, const DeclarationNameInfo &MemberNameInfo)
bool checkArgCountRange(CallExpr *Call, unsigned MinArgCount, unsigned MaxArgCount)
Checks that a call expression's argument count is in the desired range.
ExternalSemaSource * getExternalSource() const
ExprResult CreateBuiltinArraySubscriptExpr(Expr *Base, SourceLocation LLoc, Expr *Idx, SourceLocation RLoc)
bool LookupQualifiedName(LookupResult &R, DeclContext *LookupCtx, bool InUnqualifiedLookup=false)
Perform qualified name lookup into a given context.
ExprResult PerformCopyInitialization(const InitializedEntity &Entity, SourceLocation EqualLoc, ExprResult Init, bool TopLevelOfInitList=false, bool AllowExplicit=false)
ExprResult CreateBuiltinMatrixSubscriptExpr(Expr *Base, Expr *RowIdx, Expr *ColumnIdx, SourceLocation RBLoc)
Encodes a location in the source.
SourceLocation getLocWithOffset(IntTy Offset) const
Return a source location with the specified offset from this SourceLocation.
A trivial tuple used to represent a source range.
SourceLocation getEnd() const
SourceLocation getEndLoc() const LLVM_READONLY
void printPretty(raw_ostream &OS, PrinterHelper *Helper, const PrintingPolicy &Policy, unsigned Indentation=0, StringRef NewlineSymbol="\n", const ASTContext *Context=nullptr) const
SourceRange getSourceRange() const LLVM_READONLY
SourceLocation tokens are not useful in isolation - they are low level value objects created/interpre...
SourceLocation getBeginLoc() const LLVM_READONLY
StringLiteral - This represents a string literal expression, e.g.
static StringLiteral * Create(const ASTContext &Ctx, StringRef Str, StringLiteralKind Kind, bool Pascal, QualType Ty, ArrayRef< SourceLocation > Locs)
This is the "fully general" constructor that allows representation of strings formed from one or more...
void startDefinition()
Starts the definition of this tag declaration.
Exposes information about the current target.
TargetOptions & getTargetOpts() const
Retrieve the target options.
const llvm::Triple & getTriple() const
Returns the target triple of the primary target.
StringRef getPlatformName() const
Retrieve the name of the platform as it is used in the availability attribute.
VersionTuple getPlatformMinVersion() const
Retrieve the minimum desired version of the platform, to which the program should be compiled.
std::string HLSLEntry
The entry point name for HLSL shader being compiled as specified by -E.
A convenient class for passing around template argument information.
void addArgument(const TemplateArgumentLoc &Loc)
The base class of all kinds of template declarations (e.g., class, function, etc.).
Stores a list of template parameters for a TemplateDecl and its derived classes.
The top declaration context.
SourceLocation getBeginLoc() const
Get the begin source location.
A container of type source information.
TypeLoc getTypeLoc() const
Return the TypeLoc wrapper for the type source info.
The base class of the type hierarchy.
bool isBooleanType() const
bool isIncompleteArrayType() const
CXXRecordDecl * getAsCXXRecordDecl() const
Retrieves the CXXRecordDecl that this type refers to, either because the type is a RecordType or beca...
bool isConstantArrayType() const
bool hasIntegerRepresentation() const
Determine whether this type has an integer representation of some sort, e.g., it is an integer type o...
CXXRecordDecl * castAsCXXRecordDecl() const
bool isArithmeticType() const
bool isConstantMatrixType() const
bool isHLSLBuiltinIntangibleType() const
bool isPointerType() const
CanQualType getCanonicalTypeUnqualified() const
bool isIntegerType() const
isIntegerType() does not include complex integers (a GCC extension).
const T * castAs() const
Member-template castAs<specific type>.
bool isReferenceType() const
bool isHLSLIntangibleType() const
bool isEnumeralType() const
bool isScalarType() const
bool isIntegralType(const ASTContext &Ctx) const
Determine whether this type is an integral type.
const Type * getArrayElementTypeNoTypeQual() const
If this is an array type, return the element type of the array, potentially with type qualifiers miss...
QualType getPointeeType() const
If this is a pointer, ObjC object pointer, or block pointer, this returns the respective pointee.
bool hasUnsignedIntegerRepresentation() const
Determine whether this type has an unsigned integer representation of some sort, e....
bool isAggregateType() const
Determines whether the type is a C++ aggregate type or C aggregate or union type.
ScalarTypeKind getScalarTypeKind() const
Given that this is a scalar type, classify it.
bool hasSignedIntegerRepresentation() const
Determine whether this type has an signed integer representation of some sort, e.g....
bool isMatrixType() const
bool isHLSLResourceRecord() const
bool hasFloatingRepresentation() const
Determine whether this type has a floating-point representation of some sort, e.g....
bool isVectorType() const
bool isRealFloatingType() const
Floating point categories.
bool isHLSLAttributedResourceType() const
bool isFloatingType() const
const T * getAs() const
Member-template getAs<specific type>'.
const Type * getUnqualifiedDesugaredType() const
Return the specified type with any "sugar" removed from the type, removing any typedefs,...
bool isRecordType() const
bool isHLSLResourceRecordArray() const
void setType(QualType newType)
Represents a variable declaration or definition.
static VarDecl * Create(ASTContext &C, DeclContext *DC, SourceLocation StartLoc, SourceLocation IdLoc, const IdentifierInfo *Id, QualType T, TypeSourceInfo *TInfo, StorageClass S)
void setInitStyle(InitializationStyle Style)
@ CallInit
Call-style initialization (C++98)
void setStorageClass(StorageClass SC)
bool hasGlobalStorage() const
Returns true for all variables that do not have local storage.
StorageClass getStorageClass() const
Returns the storage class as written in the source.
Represents a GCC generic vector type.
unsigned getNumElements() const
QualType getElementType() const
void pushName(llvm::StringRef N)
void pushArrayIndex(uint64_t Index)
void pushBaseName(llvm::StringRef N)
IdentifierInfo * getNameAsIdentifier(ASTContext &AST) const
Defines the clang::TargetInfo interface.
uint32_t getResourceDimensions(llvm::dxil::ResourceDimension Dim)
The JSON file list parser is used to communicate input to InstallAPI.
bool isa(CodeGen::Address addr)
if(T->getSizeExpr()) TRY_TO(TraverseStmt(const_cast< Expr * >(T -> getSizeExpr())))
static bool CheckFloatOrHalfRepresentation(Sema *S, SourceLocation Loc, int ArgOrdinal, clang::QualType PassedType)
@ ICIS_NoInit
No in-class initializer.
@ TemplateName
The identifier is a template name. FIXME: Add an annotation for that.
@ OK_Ordinary
An ordinary object is located at an address in memory.
static bool CheckAllArgTypesAreCorrect(Sema *S, CallExpr *TheCall, llvm::ArrayRef< llvm::function_ref< bool(Sema *, SourceLocation, int, QualType)> > Checks)
@ AANT_ArgumentIdentifier
@ Result
The result type of a method or function.
@ Ordinary
This parameter uses ordinary ABI rules for its type.
llvm::Expected< QualType > ExpectedType
@ Template
We are parsing a template declaration.
LLVM_READONLY bool isDigit(unsigned char c)
Return true if this character is an ASCII digit: [0-9].
static bool CheckAllArgsHaveSameType(Sema *S, CallExpr *TheCall)
@ Type
The name was classified as a type.
LangAS
Defines the address space values used by the address space qualifier of QualType.
bool CreateHLSLAttributedResourceType(Sema &S, QualType Wrapped, ArrayRef< const Attr * > AttrList, QualType &ResType, HLSLAttributedResourceLocInfo *LocInfo=nullptr)
CastKind
CastKind - The kind of operation required for a conversion.
ExprValueKind
The categorization of expression values, currently following the C++11 scheme.
@ VK_PRValue
A pr-value expression (in the C++11 taxonomy) produces a temporary value.
@ VK_LValue
An l-value expression is a reference to an object with independent storage.
DynamicRecursiveASTVisitorBase< false > DynamicRecursiveASTVisitor
U cast(CodeGen::Address addr)
@ None
No keyword precedes the qualified type name.
ActionResult< Expr * > ExprResult
Visibility
Describes the different kinds of visibility that a declaration may have.
hash_code hash_value(const clang::dependencies::ModuleID &ID)
__DEVICE__ bool isnan(float __x)
__DEVICE__ _Tp abs(const std::complex< _Tp > &__c)
TypeSourceInfo * ContainedTyInfo
Describes how types, statements, expressions, and declarations should be printed.
unsigned getImplicitOrderID() const
void setCounterImplicitOrderID(unsigned Value) const
bool hasCounterImplicitOrderID() const
unsigned getSpace() const
bool hasImplicitOrderID() const
void setImplicitOrderID(unsigned Value) const
const SourceLocation & getLocation() const
const llvm::hlsl::rootsig::RootElement & getElement() const