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;
350static const HLSLAttributedResourceType *
353 "expected array of resource records");
355 while (
const ArrayType *AT = dyn_cast<ArrayType>(Ty))
357 return HLSLAttributedResourceType::findHandleTypeOnResource(Ty);
369 return RD->isEmpty();
398 Base.getType()->castAsCXXRecordDecl()))
409 assert(RD ==
nullptr &&
410 "there should be at most 1 record by a given name in a scope");
427 Name.append(NameBaseII->
getName());
434 size_t NameLength = Name.size();
443 Name.append(llvm::Twine(suffix).str());
444 II = &AST.
Idents.
get(Name, tok::TokenKind::identifier);
451 Name.truncate(NameLength);
492 "struct is already HLSL buffer compatible");
506 LS->
addAttr(PackedAttr::CreateImplicit(AST));
510 if (
unsigned NumBases = StructDecl->
getNumBases()) {
511 assert(NumBases == 1 &&
"HLSL supports only one base type");
561 LS->
addAttr(PackedAttr::CreateImplicit(AST));
566 VarDecl *VD = dyn_cast<VarDecl>(D);
586 uint32_t ImplicitBindingOrderID) {
588 HLSLResourceBindingAttr::CreateImplicit(S.
getASTContext(),
"",
"0", {});
589 Attr->setBinding(RT, std::nullopt, 0);
590 Attr->setImplicitBindingOrderID(ImplicitBindingOrderID);
597 BufDecl->setRBraceLoc(RBrace);
609 uint32_t OrderID = getNextImplicitBindingOrderID();
614 BufDecl->isCBuffer() ? RegisterType::CBuffer
624 int X,
int Y,
int Z) {
625 if (HLSLNumThreadsAttr *NT = D->
getAttr<HLSLNumThreadsAttr>()) {
626 if (NT->getX() !=
X || NT->getY() != Y || NT->getZ() != Z) {
627 Diag(NT->getLocation(), diag::err_hlsl_attribute_param_mismatch) << AL;
628 Diag(AL.
getLoc(), diag::note_conflicting_attribute);
638 int Min,
int Max,
int Preferred,
639 int SpelledArgsCount) {
640 if (HLSLWaveSizeAttr *WS = D->
getAttr<HLSLWaveSizeAttr>()) {
641 if (WS->getMin() !=
Min || WS->getMax() !=
Max ||
642 WS->getPreferred() != Preferred ||
643 WS->getSpelledArgsCount() != SpelledArgsCount) {
644 Diag(WS->getLocation(), diag::err_hlsl_attribute_param_mismatch) << AL;
645 Diag(AL.
getLoc(), diag::note_conflicting_attribute);
651 Result->setSpelledArgsCount(SpelledArgsCount);
655HLSLVkConstantIdAttr *
661 Diag(AL.
getLoc(), diag::warn_attribute_ignored) << AL;
669 Diag(VD->getLocation(), diag::err_specialization_const);
673 if (!VD->getType().isConstQualified()) {
674 Diag(VD->getLocation(), diag::err_specialization_const);
678 if (HLSLVkConstantIdAttr *CI = D->
getAttr<HLSLVkConstantIdAttr>()) {
679 if (CI->getId() != Id) {
680 Diag(CI->getLocation(), diag::err_hlsl_attribute_param_mismatch) << AL;
681 Diag(AL.
getLoc(), diag::note_conflicting_attribute);
686 HLSLVkConstantIdAttr *
Result =
693 llvm::Triple::EnvironmentType ShaderType) {
694 if (HLSLShaderAttr *NT = D->
getAttr<HLSLShaderAttr>()) {
695 if (NT->getType() != ShaderType) {
696 Diag(NT->getLocation(), diag::err_hlsl_attribute_param_mismatch) << AL;
697 Diag(AL.
getLoc(), diag::note_conflicting_attribute);
701 return HLSLShaderAttr::Create(
getASTContext(), ShaderType, AL);
704HLSLParamModifierAttr *
706 HLSLParamModifierAttr::Spelling Spelling) {
709 if (HLSLParamModifierAttr *PA = D->
getAttr<HLSLParamModifierAttr>()) {
710 if ((PA->isIn() && Spelling == HLSLParamModifierAttr::Keyword_out) ||
711 (PA->isOut() && Spelling == HLSLParamModifierAttr::Keyword_in)) {
712 D->
dropAttr<HLSLParamModifierAttr>();
714 return HLSLParamModifierAttr::Create(
716 HLSLParamModifierAttr::Keyword_inout);
718 Diag(AL.
getLoc(), diag::err_hlsl_duplicate_parameter_modifier) << AL;
719 Diag(PA->getLocation(), diag::note_conflicting_attribute);
745 if (HLSLShaderAttr::isValidShaderType(Env) && Env != llvm::Triple::Library) {
746 if (
const auto *Shader = FD->
getAttr<HLSLShaderAttr>()) {
749 if (Shader->getType() != Env) {
750 Diag(Shader->getLocation(), diag::err_hlsl_entry_shader_attr_mismatch)
762 case llvm::Triple::UnknownEnvironment:
763 case llvm::Triple::Library:
765 case llvm::Triple::RootSignature:
766 llvm_unreachable(
"rootsig environment has no functions");
768 llvm_unreachable(
"Unhandled environment in triple");
774 HLSLAppliedSemanticAttr *Semantic,
779 const auto *ShaderAttr = FD->
getAttr<HLSLShaderAttr>();
780 assert(ShaderAttr &&
"Entry point has no shader attribute");
781 llvm::Triple::EnvironmentType ST = ShaderAttr->getType();
782 auto SemanticName = Semantic->getSemanticName().upper();
787 if (SemanticName ==
"SV_POSITION") {
788 return (ST == llvm::Triple::Vertex && !IsInput) ||
789 (ST == llvm::Triple::Pixel && IsInput);
795bool SemaHLSL::determineActiveSemanticOnScalar(
FunctionDecl *FD,
798 SemanticInfo &ActiveSemantic,
799 SemaHLSL::SemanticContext &SC) {
800 if (ActiveSemantic.Semantic ==
nullptr) {
801 ActiveSemantic.Semantic = D->
getAttr<HLSLParsedSemanticAttr>();
802 if (ActiveSemantic.Semantic)
803 ActiveSemantic.Index = ActiveSemantic.Semantic->getSemanticIndex();
806 if (!ActiveSemantic.Semantic) {
812 HLSLAppliedSemanticAttr(
getASTContext(), *ActiveSemantic.Semantic,
813 ActiveSemantic.Semantic->getAttrName()->getName(),
814 ActiveSemantic.Index.value_or(0));
818 checkSemanticAnnotation(FD, D, A, SC);
819 OutputDecl->addAttr(A);
821 unsigned Location = ActiveSemantic.Index.value_or(0);
824 SC.CurrentIOType & IOType::In)) {
825 bool HasVkLocation =
false;
826 if (
auto *A = D->getAttr<HLSLVkLocationAttr>()) {
827 HasVkLocation = true;
828 Location = A->getLocation();
831 if (SC.UsesExplicitVkLocations.value_or(HasVkLocation) != HasVkLocation) {
832 Diag(D->getLocation(), diag::err_hlsl_semantic_partial_explicit_indexing);
835 SC.UsesExplicitVkLocations = HasVkLocation;
838 const ConstantArrayType *AT = dyn_cast<ConstantArrayType>(D->getType());
839 unsigned ElementCount = AT ? AT->
getZExtSize() : 1;
840 ActiveSemantic.Index = Location + ElementCount;
842 Twine BaseName = Twine(ActiveSemantic.Semantic->getAttrName()->getName());
843 for (
unsigned I = 0; I < ElementCount; ++I) {
844 Twine VariableName = BaseName.concat(Twine(Location + I));
846 auto [_, Inserted] = SC.ActiveSemantics.insert(VariableName.str());
848 Diag(D->getLocation(), diag::err_hlsl_semantic_index_overlap)
849 << VariableName.str();
860 SemanticInfo &ActiveSemantic,
861 SemaHLSL::SemanticContext &SC) {
862 if (ActiveSemantic.Semantic ==
nullptr) {
863 ActiveSemantic.Semantic = D->
getAttr<HLSLParsedSemanticAttr>();
864 if (ActiveSemantic.Semantic)
865 ActiveSemantic.Index = ActiveSemantic.Semantic->getSemanticIndex();
871 const RecordType *RT = dyn_cast<RecordType>(T);
873 return determineActiveSemanticOnScalar(FD, OutputDecl, D, ActiveSemantic,
876 const RecordDecl *RD = RT->getDecl();
877 for (FieldDecl *Field : RD->
fields()) {
878 SemanticInfo Info = ActiveSemantic;
879 if (!determineActiveSemantic(FD, OutputDecl, Field, Info, SC)) {
880 Diag(
Field->getLocation(), diag::note_hlsl_semantic_used_here) <<
Field;
883 if (ActiveSemantic.Semantic)
884 ActiveSemantic = Info;
891 const auto *ShaderAttr = FD->
getAttr<HLSLShaderAttr>();
892 assert(ShaderAttr &&
"Entry point has no shader attribute");
893 llvm::Triple::EnvironmentType ST = ShaderAttr->getType();
897 case llvm::Triple::Pixel:
898 case llvm::Triple::Vertex:
899 case llvm::Triple::Geometry:
900 case llvm::Triple::Hull:
901 case llvm::Triple::Domain:
902 case llvm::Triple::RayGeneration:
903 case llvm::Triple::Intersection:
904 case llvm::Triple::AnyHit:
905 case llvm::Triple::ClosestHit:
906 case llvm::Triple::Miss:
907 case llvm::Triple::Callable:
908 if (
const auto *NT = FD->
getAttr<HLSLNumThreadsAttr>()) {
909 diagnoseAttrStageMismatch(NT, ST,
910 {llvm::Triple::Compute,
911 llvm::Triple::Amplification,
912 llvm::Triple::Mesh});
915 if (
const auto *WS = FD->
getAttr<HLSLWaveSizeAttr>()) {
916 diagnoseAttrStageMismatch(WS, ST,
917 {llvm::Triple::Compute,
918 llvm::Triple::Amplification,
919 llvm::Triple::Mesh});
924 case llvm::Triple::Compute:
925 case llvm::Triple::Amplification:
926 case llvm::Triple::Mesh:
927 if (!FD->
hasAttr<HLSLNumThreadsAttr>()) {
929 << llvm::Triple::getEnvironmentTypeName(ST);
932 if (
const auto *WS = FD->
getAttr<HLSLWaveSizeAttr>()) {
933 if (Ver < VersionTuple(6, 6)) {
934 Diag(WS->getLocation(), diag::err_hlsl_attribute_in_wrong_shader_model)
937 }
else if (WS->getSpelledArgsCount() > 1 && Ver < VersionTuple(6, 8)) {
940 diag::err_hlsl_attribute_number_arguments_insufficient_shader_model)
941 << WS << WS->getSpelledArgsCount() <<
"6.8";
946 case llvm::Triple::RootSignature:
947 llvm_unreachable(
"rootsig environment has no function entry point");
949 llvm_unreachable(
"Unhandled environment in triple");
952 SemaHLSL::SemanticContext InputSC = {};
953 InputSC.CurrentIOType = IOType::In;
956 SemanticInfo ActiveSemantic;
957 ActiveSemantic.Semantic = Param->getAttr<HLSLParsedSemanticAttr>();
958 if (ActiveSemantic.Semantic)
959 ActiveSemantic.Index = ActiveSemantic.Semantic->getSemanticIndex();
962 if (!determineActiveSemantic(FD, Param, Param, ActiveSemantic, InputSC)) {
963 Diag(Param->getLocation(), diag::note_previous_decl) << Param;
968 SemanticInfo ActiveSemantic;
969 SemaHLSL::SemanticContext OutputSC = {};
970 OutputSC.CurrentIOType = IOType::Out;
971 ActiveSemantic.Semantic = FD->
getAttr<HLSLParsedSemanticAttr>();
972 if (ActiveSemantic.Semantic)
973 ActiveSemantic.Index = ActiveSemantic.Semantic->getSemanticIndex();
975 determineActiveSemantic(FD, FD, FD, ActiveSemantic, OutputSC);
978void SemaHLSL::checkSemanticAnnotation(
980 const HLSLAppliedSemanticAttr *SemanticAttr,
const SemanticContext &SC) {
981 auto *ShaderAttr = EntryPoint->
getAttr<HLSLShaderAttr>();
982 assert(ShaderAttr &&
"Entry point has no shader attribute");
983 llvm::Triple::EnvironmentType ST = ShaderAttr->getType();
985 auto SemanticName = SemanticAttr->getSemanticName().upper();
986 if (SemanticName ==
"SV_DISPATCHTHREADID" ||
987 SemanticName ==
"SV_GROUPINDEX" || SemanticName ==
"SV_GROUPTHREADID" ||
988 SemanticName ==
"SV_GROUPID") {
990 if (ST != llvm::Triple::Compute)
991 diagnoseSemanticStageMismatch(SemanticAttr, ST, SC.CurrentIOType,
992 {{llvm::Triple::Compute, IOType::In}});
994 if (SemanticAttr->getSemanticIndex() != 0) {
995 std::string PrettyName =
996 "'" + SemanticAttr->getSemanticName().str() +
"'";
997 Diag(SemanticAttr->getLoc(),
998 diag::err_hlsl_semantic_indexing_not_supported)
1004 if (SemanticName ==
"SV_POSITION") {
1007 diagnoseSemanticStageMismatch(SemanticAttr, ST, SC.CurrentIOType,
1008 {{llvm::Triple::Vertex, IOType::InOut},
1009 {llvm::Triple::Pixel, IOType::In}});
1013 if (SemanticName ==
"SV_TARGET") {
1014 diagnoseSemanticStageMismatch(SemanticAttr, ST, SC.CurrentIOType,
1015 {{llvm::Triple::Pixel, IOType::Out}});
1021 if (SemanticAttr->getAttrName()->getName().starts_with_insensitive(
"SV_"))
1022 llvm_unreachable(
"Unknown SemanticAttr");
1025void SemaHLSL::diagnoseAttrStageMismatch(
1026 const Attr *A, llvm::Triple::EnvironmentType Stage,
1027 std::initializer_list<llvm::Triple::EnvironmentType> AllowedStages) {
1028 SmallVector<StringRef, 8> StageStrings;
1029 llvm::transform(AllowedStages, std::back_inserter(StageStrings),
1030 [](llvm::Triple::EnvironmentType ST) {
1032 HLSLShaderAttr::ConvertEnvironmentTypeToStr(ST));
1034 Diag(A->
getLoc(), diag::err_hlsl_attr_unsupported_in_stage)
1035 << A->
getAttrName() << llvm::Triple::getEnvironmentTypeName(Stage)
1036 << (AllowedStages.size() != 1) << join(StageStrings,
", ");
1039void SemaHLSL::diagnoseSemanticStageMismatch(
1040 const Attr *A, llvm::Triple::EnvironmentType Stage, IOType CurrentIOType,
1041 std::initializer_list<SemanticStageInfo> Allowed) {
1043 for (
auto &Case : Allowed) {
1044 if (Case.Stage != Stage)
1047 if (CurrentIOType & Case.AllowedIOTypesMask)
1050 SmallVector<std::string, 8> ValidCases;
1052 Allowed, std::back_inserter(ValidCases), [](SemanticStageInfo Case) {
1053 SmallVector<std::string, 2> ValidType;
1054 if (Case.AllowedIOTypesMask & IOType::In)
1055 ValidType.push_back(
"input");
1056 if (Case.AllowedIOTypesMask & IOType::Out)
1057 ValidType.push_back(
"output");
1059 HLSLShaderAttr::ConvertEnvironmentTypeToStr(Case.Stage)) +
1060 " " + join(ValidType,
"/");
1062 Diag(A->
getLoc(), diag::err_hlsl_semantic_unsupported_iotype_for_stage)
1063 << A->
getAttrName() << (CurrentIOType & IOType::In ?
"input" :
"output")
1064 << llvm::Triple::getEnvironmentTypeName(Case.Stage)
1065 << join(ValidCases,
", ");
1069 SmallVector<StringRef, 8> StageStrings;
1071 Allowed, std::back_inserter(StageStrings), [](SemanticStageInfo Case) {
1073 HLSLShaderAttr::ConvertEnvironmentTypeToStr(Case.Stage));
1076 Diag(A->
getLoc(), diag::err_hlsl_attr_unsupported_in_stage)
1077 << A->
getAttrName() << llvm::Triple::getEnvironmentTypeName(Stage)
1078 << (Allowed.size() != 1) << join(StageStrings,
", ");
1081template <CastKind Kind>
1084 Ty = VTy->getElementType();
1089template <CastKind Kind>
1101 if (LHSFloat && RHSFloat) {
1129 if (LHSSigned == RHSSigned) {
1130 if (IsCompAssign || IntOrder >= 0)
1138 if (IntOrder != (LHSSigned ? 1 : -1)) {
1139 if (IsCompAssign || RHSSigned)
1147 if (Ctx.getIntWidth(LElTy) != Ctx.getIntWidth(RElTy)) {
1148 if (IsCompAssign || LHSSigned)
1164 QualType ElTy = Ctx.getCorrespondingUnsignedType(LHSSigned ? LElTy : RElTy);
1165 QualType NewTy = Ctx.getExtVectorType(
1175 return CK_FloatingCast;
1177 return CK_IntegralCast;
1179 return CK_IntegralToFloating;
1181 return CK_FloatingToIntegral;
1187 bool IsCompAssign) {
1194 if (!LVecTy && IsCompAssign) {
1196 RHS =
SemaRef.ImpCastExprToType(RHS.
get(), RElTy, CK_HLSLVectorTruncation);
1198 if (Ctx.hasSameUnqualifiedType(LHSType, RHSType))
1200 RHS =
SemaRef.ImpCastExprToType(RHS.
get(), LHSType,
1205 unsigned EndSz = std::numeric_limits<unsigned>::max();
1208 LSz = EndSz = LVecTy->getNumElements();
1211 assert(EndSz != std::numeric_limits<unsigned>::max() &&
1212 "one of the above should have had a value");
1216 if (IsCompAssign && LSz != EndSz) {
1218 diag::err_hlsl_vector_compound_assignment_truncation)
1219 << LHSType << RHSType;
1225 if (!IsCompAssign && LVecTy && LVecTy->getNumElements() > EndSz)
1230 if (!IsCompAssign && !LVecTy)
1234 if (Ctx.hasSameUnqualifiedType(LHSType, RHSType))
1235 return Ctx.getCommonSugaredType(LHSType, RHSType);
1243 LElTy, RElTy, IsCompAssign);
1246 "HLSL Vectors can only contain integer or floating point types");
1248 LElTy, RElTy, IsCompAssign);
1253 assert((Opc == BO_LOr || Opc == BO_LAnd) &&
1254 "Called with non-logical operator");
1256 llvm::raw_svector_ostream OS(Buff);
1258 StringRef NewFnName = Opc == BO_LOr ?
"or" :
"and";
1259 OS << NewFnName <<
"(";
1269std::pair<IdentifierInfo *, bool>
1272 std::string IdStr =
"__hlsl_rootsig_decl_" + std::to_string(Hash);
1279 return {DeclIdent,
Found};
1290 for (
auto &RootSigElement : RootElements)
1291 Elements.push_back(RootSigElement.getElement());
1295 DeclIdent,
SemaRef.getLangOpts().HLSLRootSigVer, Elements);
1297 SignatureDecl->setImplicit();
1303 if (RootSigOverrideIdent) {
1306 if (
SemaRef.LookupQualifiedName(R, DC))
1307 return dyn_cast<HLSLRootSignatureDecl>(R.
getFoundDecl());
1315struct PerVisibilityBindingChecker {
1318 std::array<llvm::hlsl::BindingInfoBuilder, 8> Builders;
1322 llvm::dxbc::ShaderVisibility Vis;
1327 PerVisibilityBindingChecker(
SemaHLSL *S) : S(S) {}
1329 void trackBinding(llvm::dxbc::ShaderVisibility
Visibility,
1330 llvm::dxil::ResourceClass RC, uint32_t Space,
1331 uint32_t LowerBound, uint32_t UpperBound,
1332 const hlsl::RootSignatureElement *Elem) {
1334 assert(BuilderIndex < Builders.size() &&
1335 "Not enough builders for visibility type");
1336 Builders[BuilderIndex].trackBinding(RC, Space, LowerBound, UpperBound,
1337 static_cast<const void *
>(Elem));
1339 static_assert(llvm::to_underlying(llvm::dxbc::ShaderVisibility::All) == 0,
1340 "'All' visibility must come first");
1341 if (
Visibility == llvm::dxbc::ShaderVisibility::All)
1342 for (
size_t I = 1, E = Builders.size(); I < E; ++I)
1343 Builders[I].trackBinding(RC, Space, LowerBound, UpperBound,
1344 static_cast<const void *
>(Elem));
1346 ElemInfoMap.push_back({Elem,
Visibility,
false});
1349 ElemInfo &
getInfo(
const hlsl::RootSignatureElement *Elem) {
1350 auto It = llvm::lower_bound(
1352 [](
const auto &LHS,
const auto &RHS) {
return LHS.Elem < RHS; });
1353 assert(It->Elem == Elem &&
"Element not in map");
1357 bool checkOverlap() {
1358 llvm::sort(ElemInfoMap, [](
const auto &LHS,
const auto &RHS) {
1359 return LHS.Elem < RHS.Elem;
1362 bool HadOverlap =
false;
1364 using llvm::hlsl::BindingInfoBuilder;
1365 auto ReportOverlap = [
this,
1366 &HadOverlap](
const BindingInfoBuilder &Builder,
1367 const llvm::hlsl::Binding &Reported) {
1371 static_cast<const hlsl::RootSignatureElement *
>(Reported.Cookie);
1372 const llvm::hlsl::Binding &
Previous = Builder.findOverlapping(Reported);
1373 const auto *PrevElem =
1374 static_cast<const hlsl::RootSignatureElement *
>(
Previous.Cookie);
1376 ElemInfo &Info =
getInfo(Elem);
1381 Info.Diagnosed =
true;
1383 ElemInfo &PrevInfo =
getInfo(PrevElem);
1384 llvm::dxbc::ShaderVisibility CommonVis =
1385 Info.Vis == llvm::dxbc::ShaderVisibility::All ? PrevInfo.Vis
1388 this->S->
Diag(Elem->
getLocation(), diag::err_hlsl_resource_range_overlap)
1389 << llvm::to_underlying(Reported.RC) << Reported.LowerBound
1390 << Reported.isUnbounded() << Reported.UpperBound
1395 this->S->
Diag(PrevElem->getLocation(),
1396 diag::note_hlsl_resource_range_here);
1399 for (BindingInfoBuilder &Builder : Builders)
1400 Builder.calculateBindingInfo(ReportOverlap);
1424 if (
const auto *ResTy =
1425 SecondField->
getType()->
getAs<HLSLAttributedResourceType>()) {
1426 return ResTy->getAttrs().IsCounter;
1434 bool HadError =
false;
1435 auto ReportError = [
this, &HadError](
SourceLocation Loc, uint32_t LowerBound,
1436 uint32_t UpperBound) {
1438 this->
Diag(Loc, diag::err_hlsl_invalid_rootsig_value)
1439 << LowerBound << UpperBound;
1446 this->
Diag(Loc, diag::err_hlsl_invalid_rootsig_value)
1447 << llvm::formatv(
"{0:f}", LowerBound).sstr<6>()
1448 << llvm::formatv(
"{0:f}", UpperBound).sstr<6>();
1451 auto VerifyRegister = [ReportError](
SourceLocation Loc, uint32_t Register) {
1452 if (!llvm::hlsl::rootsig::verifyRegisterValue(Register))
1453 ReportError(Loc, 0, 0xfffffffe);
1456 auto VerifySpace = [ReportError](
SourceLocation Loc, uint32_t Space) {
1457 if (!llvm::hlsl::rootsig::verifyRegisterSpace(Space))
1458 ReportError(Loc, 0, 0xffffffef);
1461 const uint32_t Version =
1462 llvm::to_underlying(
SemaRef.getLangOpts().HLSLRootSigVer);
1463 const uint32_t VersionEnum = Version - 1;
1464 auto ReportFlagError = [
this, &HadError, VersionEnum](
SourceLocation Loc) {
1466 this->
Diag(Loc, diag::err_hlsl_invalid_rootsig_flag)
1473 const llvm::hlsl::rootsig::RootElement &Elem = RootSigElem.
getElement();
1474 if (
const auto *Descriptor =
1475 std::get_if<llvm::hlsl::rootsig::RootDescriptor>(&Elem)) {
1476 VerifyRegister(Loc, Descriptor->Reg.Number);
1477 VerifySpace(Loc, Descriptor->Space);
1479 if (!llvm::hlsl::rootsig::verifyRootDescriptorFlag(Version,
1481 ReportFlagError(Loc);
1482 }
else if (
const auto *Constants =
1483 std::get_if<llvm::hlsl::rootsig::RootConstants>(&Elem)) {
1484 VerifyRegister(Loc, Constants->Reg.Number);
1485 VerifySpace(Loc, Constants->Space);
1486 }
else if (
const auto *Sampler =
1487 std::get_if<llvm::hlsl::rootsig::StaticSampler>(&Elem)) {
1488 VerifyRegister(Loc, Sampler->Reg.Number);
1489 VerifySpace(Loc, Sampler->Space);
1492 "By construction, parseFloatParam can't produce a NaN from a "
1493 "float_literal token");
1495 if (!llvm::hlsl::rootsig::verifyMaxAnisotropy(Sampler->MaxAnisotropy))
1496 ReportError(Loc, 0, 16);
1497 if (!llvm::hlsl::rootsig::verifyMipLODBias(Sampler->MipLODBias))
1498 ReportFloatError(Loc, -16.f, 15.99f);
1499 }
else if (
const auto *Clause =
1500 std::get_if<llvm::hlsl::rootsig::DescriptorTableClause>(
1502 VerifyRegister(Loc, Clause->Reg.Number);
1503 VerifySpace(Loc, Clause->Space);
1505 if (!llvm::hlsl::rootsig::verifyNumDescriptors(Clause->NumDescriptors)) {
1509 ReportError(Loc, 1, 0xfffffffe);
1512 if (!llvm::hlsl::rootsig::verifyDescriptorRangeFlag(Version, Clause->Type,
1514 ReportFlagError(Loc);
1518 PerVisibilityBindingChecker BindingChecker(
this);
1519 SmallVector<std::pair<
const llvm::hlsl::rootsig::DescriptorTableClause *,
1524 const llvm::hlsl::rootsig::RootElement &Elem = RootSigElem.
getElement();
1525 if (
const auto *Descriptor =
1526 std::get_if<llvm::hlsl::rootsig::RootDescriptor>(&Elem)) {
1527 uint32_t LowerBound(Descriptor->Reg.Number);
1528 uint32_t UpperBound(LowerBound);
1530 BindingChecker.trackBinding(
1531 Descriptor->Visibility,
1532 static_cast<llvm::dxil::ResourceClass
>(Descriptor->Type),
1533 Descriptor->Space, LowerBound, UpperBound, &RootSigElem);
1534 }
else if (
const auto *Constants =
1535 std::get_if<llvm::hlsl::rootsig::RootConstants>(&Elem)) {
1536 uint32_t LowerBound(Constants->Reg.Number);
1537 uint32_t UpperBound(LowerBound);
1539 BindingChecker.trackBinding(
1540 Constants->Visibility, llvm::dxil::ResourceClass::CBuffer,
1541 Constants->Space, LowerBound, UpperBound, &RootSigElem);
1542 }
else if (
const auto *Sampler =
1543 std::get_if<llvm::hlsl::rootsig::StaticSampler>(&Elem)) {
1544 uint32_t LowerBound(Sampler->Reg.Number);
1545 uint32_t UpperBound(LowerBound);
1547 BindingChecker.trackBinding(
1548 Sampler->Visibility, llvm::dxil::ResourceClass::Sampler,
1549 Sampler->Space, LowerBound, UpperBound, &RootSigElem);
1550 }
else if (
const auto *Clause =
1551 std::get_if<llvm::hlsl::rootsig::DescriptorTableClause>(
1554 UnboundClauses.emplace_back(Clause, &RootSigElem);
1555 }
else if (
const auto *Table =
1556 std::get_if<llvm::hlsl::rootsig::DescriptorTable>(&Elem)) {
1557 assert(UnboundClauses.size() == Table->NumClauses &&
1558 "Number of unbound elements must match the number of clauses");
1559 bool HasAnySampler =
false;
1560 bool HasAnyNonSampler =
false;
1561 uint64_t Offset = 0;
1562 bool IsPrevUnbound =
false;
1563 for (
const auto &[Clause, ClauseElem] : UnboundClauses) {
1565 if (Clause->Type == llvm::dxil::ResourceClass::Sampler)
1566 HasAnySampler =
true;
1568 HasAnyNonSampler =
true;
1570 if (HasAnySampler && HasAnyNonSampler)
1571 Diag(Loc, diag::err_hlsl_invalid_mixed_resources);
1576 if (Clause->NumDescriptors == 0)
1580 Clause->Offset == llvm::hlsl::rootsig::DescriptorTableOffsetAppend;
1582 Offset = Clause->Offset;
1584 uint64_t RangeBound = llvm::hlsl::rootsig::computeRangeBound(
1585 Offset, Clause->NumDescriptors);
1587 if (IsPrevUnbound && IsAppending)
1588 Diag(Loc, diag::err_hlsl_appending_onto_unbound);
1589 else if (!llvm::hlsl::rootsig::verifyNoOverflowedOffset(RangeBound))
1590 Diag(Loc, diag::err_hlsl_offset_overflow) << Offset << RangeBound;
1593 Offset = RangeBound + 1;
1594 IsPrevUnbound = Clause->NumDescriptors ==
1595 llvm::hlsl::rootsig::NumDescriptorsUnbounded;
1598 uint32_t LowerBound(Clause->Reg.Number);
1599 uint32_t UpperBound = llvm::hlsl::rootsig::computeRangeBound(
1600 LowerBound, Clause->NumDescriptors);
1602 BindingChecker.trackBinding(
1604 static_cast<llvm::dxil::ResourceClass
>(Clause->Type), Clause->Space,
1605 LowerBound, UpperBound, ClauseElem);
1607 UnboundClauses.clear();
1611 return BindingChecker.checkOverlap();
1616 Diag(AL.
getLoc(), diag::err_attribute_wrong_number_arguments) << AL << 1;
1621 if (
auto *RS = D->
getAttr<RootSignatureAttr>()) {
1622 if (RS->getSignatureIdent() != Ident) {
1623 Diag(AL.
getLoc(), diag::err_disallowed_duplicate_attribute) << RS;
1627 Diag(AL.
getLoc(), diag::warn_duplicate_attribute_exact) << RS;
1633 if (
auto *SignatureDecl =
1641 llvm::VersionTuple SMVersion =
1646 uint32_t ZMax = 1024;
1647 uint32_t ThreadMax = 1024;
1648 if (IsDXIL && SMVersion.getMajor() <= 4) {
1651 }
else if (IsDXIL && SMVersion.getMajor() == 5) {
1661 diag::err_hlsl_numthreads_argument_oor)
1670 diag::err_hlsl_numthreads_argument_oor)
1679 diag::err_hlsl_numthreads_argument_oor)
1684 if (
X * Y * Z > ThreadMax) {
1685 Diag(AL.
getLoc(), diag::err_hlsl_numthreads_invalid) << ThreadMax;
1702 if (SpelledArgsCount == 0 || SpelledArgsCount > 3)
1710 if (SpelledArgsCount > 1 &&
1714 uint32_t Preferred = 0;
1715 if (SpelledArgsCount > 2 &&
1719 if (SpelledArgsCount > 2) {
1722 diag::err_attribute_power_of_two_in_range)
1723 << AL << llvm::dxil::MinWaveSize << llvm::dxil::MaxWaveSize
1728 if (Preferred < Min || Preferred >
Max) {
1730 diag::err_attribute_power_of_two_in_range)
1731 << AL <<
Min <<
Max << Preferred;
1734 }
else if (SpelledArgsCount > 1) {
1737 diag::err_attribute_power_of_two_in_range)
1738 << AL << llvm::dxil::MinWaveSize << llvm::dxil::MaxWaveSize <<
Max;
1742 Diag(AL.
getLoc(), diag::err_attribute_argument_invalid) << AL << 1;
1745 Diag(AL.
getLoc(), diag::warn_attr_min_eq_max) << AL;
1750 diag::err_attribute_power_of_two_in_range)
1751 << AL << llvm::dxil::MinWaveSize << llvm::dxil::MaxWaveSize <<
Min;
1756 HLSLWaveSizeAttr *NewAttr =
1785 uint32_t Binding = 0;
1809 if (!T->hasUnsignedIntegerRepresentation() ||
1810 (VT && VT->getNumElements() > 3)) {
1811 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type)
1812 << AL <<
"uint/uint2/uint3";
1821 if (!T->hasFloatingRepresentation() || (VT && VT->getNumElements() > 4)) {
1822 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type)
1823 << AL <<
"float/float1/float2/float3/float4";
1831 std::optional<unsigned> Index) {
1835 QualType ValueType = VD->getType();
1836 if (
auto *FD = dyn_cast<FunctionDecl>(D))
1839 bool IsOutput =
false;
1840 if (HLSLParamModifierAttr *MA = D->
getAttr<HLSLParamModifierAttr>()) {
1847 if (SemanticName ==
"SV_DISPATCHTHREADID") {
1850 Diag(AL.
getLoc(), diag::err_hlsl_semantic_output_not_supported) << AL;
1851 if (Index.has_value())
1852 Diag(AL.
getLoc(), diag::err_hlsl_semantic_indexing_not_supported) << AL;
1857 if (SemanticName ==
"SV_GROUPINDEX") {
1859 Diag(AL.
getLoc(), diag::err_hlsl_semantic_output_not_supported) << AL;
1860 if (Index.has_value())
1861 Diag(AL.
getLoc(), diag::err_hlsl_semantic_indexing_not_supported) << AL;
1866 if (SemanticName ==
"SV_GROUPTHREADID") {
1869 Diag(AL.
getLoc(), diag::err_hlsl_semantic_output_not_supported) << AL;
1870 if (Index.has_value())
1871 Diag(AL.
getLoc(), diag::err_hlsl_semantic_indexing_not_supported) << AL;
1876 if (SemanticName ==
"SV_GROUPID") {
1879 Diag(AL.
getLoc(), diag::err_hlsl_semantic_output_not_supported) << AL;
1880 if (Index.has_value())
1881 Diag(AL.
getLoc(), diag::err_hlsl_semantic_indexing_not_supported) << AL;
1886 if (SemanticName ==
"SV_POSITION") {
1887 const auto *VT = ValueType->getAs<
VectorType>();
1888 if (!ValueType->hasFloatingRepresentation() ||
1889 (VT && VT->getNumElements() > 4))
1890 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type)
1891 << AL <<
"float/float1/float2/float3/float4";
1896 if (SemanticName ==
"SV_TARGET") {
1897 const auto *VT = ValueType->getAs<
VectorType>();
1898 if (!ValueType->hasFloatingRepresentation() ||
1899 (VT && VT->getNumElements() > 4))
1900 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_type)
1901 << AL <<
"float/float1/float2/float3/float4";
1906 Diag(AL.
getLoc(), diag::err_hlsl_unknown_semantic) << AL;
1910 uint32_t IndexValue(0), ExplicitIndex(0);
1913 assert(0 &&
"HLSLUnparsedSemantic is expected to have 2 int arguments.");
1915 assert(IndexValue > 0 ? ExplicitIndex :
true);
1916 std::optional<unsigned> Index =
1917 ExplicitIndex ? std::optional<unsigned>(IndexValue) : std::nullopt;
1927 Diag(AL.
getLoc(), diag::err_hlsl_attr_invalid_ast_node)
1928 << AL <<
"shader constant in a constant buffer";
1932 uint32_t SubComponent;
1942 bool IsAggregateTy = (T->isArrayType() || T->isStructureType());
1947 if (IsAggregateTy) {
1948 Diag(AL.
getLoc(), diag::err_hlsl_invalid_register_or_packoffset);
1952 if ((Component * 32 + Size) > 128) {
1953 Diag(AL.
getLoc(), diag::err_hlsl_packoffset_cross_reg_boundary);
1958 EltTy = VT->getElementType();
1960 if (Align > 32 && Component == 1) {
1963 Diag(AL.
getLoc(), diag::err_hlsl_packoffset_alignment_mismatch)
1977 if (!
SemaRef.checkStringLiteralArgumentAttr(AL, 0, Str, &ArgLoc))
1980 llvm::Triple::EnvironmentType ShaderType;
1981 if (!HLSLShaderAttr::ConvertStrToEnvironmentType(Str, ShaderType)) {
1982 Diag(AL.
getLoc(), diag::warn_attribute_type_not_supported)
1983 << AL << Str << ArgLoc;
1997 assert(AttrList.size() &&
"expected list of resource attributes");
2004 HLSLAttributedResourceType::Attributes ResAttrs;
2006 bool HasResourceClass =
false;
2007 bool HasResourceDimension =
false;
2008 for (
const Attr *A : AttrList) {
2013 case attr::HLSLResourceClass: {
2015 if (HasResourceClass) {
2017 ? diag::warn_duplicate_attribute_exact
2018 : diag::warn_duplicate_attribute)
2022 ResAttrs.ResourceClass = RC;
2023 HasResourceClass =
true;
2026 case attr::HLSLResourceDimension: {
2027 llvm::dxil::ResourceDimension RD =
2029 if (HasResourceDimension) {
2031 ? diag::warn_duplicate_attribute_exact
2032 : diag::warn_duplicate_attribute)
2036 ResAttrs.ResourceDimension = RD;
2037 HasResourceDimension =
true;
2041 if (ResAttrs.IsROV) {
2045 ResAttrs.IsROV =
true;
2047 case attr::HLSLRawBuffer:
2048 if (ResAttrs.RawBuffer) {
2052 ResAttrs.RawBuffer =
true;
2054 case attr::HLSLIsCounter:
2055 if (ResAttrs.IsCounter) {
2059 ResAttrs.IsCounter =
true;
2061 case attr::HLSLContainedType: {
2064 if (!ContainedTy.
isNull()) {
2066 ? diag::warn_duplicate_attribute_exact
2067 : diag::warn_duplicate_attribute)
2076 llvm_unreachable(
"unhandled resource attribute type");
2080 if (!HasResourceClass) {
2081 S.
Diag(AttrList.back()->getRange().getEnd(),
2082 diag::err_hlsl_missing_resource_class);
2087 Wrapped, ContainedTy, ResAttrs);
2089 if (LocInfo && ContainedTyInfo) {
2102 if (!T->isHLSLResourceType()) {
2103 Diag(AL.
getLoc(), diag::err_hlsl_attribute_needs_intangible_type)
2118 AttributeCommonInfo::AS_CXX11, 0, false ,
2123 case ParsedAttr::AT_HLSLResourceClass: {
2125 Diag(AL.
getLoc(), diag::err_attribute_argument_type)
2136 if (!HLSLResourceClassAttr::ConvertStrToResourceClass(Identifier, RC)) {
2137 Diag(ArgLoc, diag::warn_attribute_type_not_supported)
2138 <<
"ResourceClass" << Identifier;
2141 A = HLSLResourceClassAttr::Create(
getASTContext(), RC, ACI);
2145 case ParsedAttr::AT_HLSLROV:
2149 case ParsedAttr::AT_HLSLRawBuffer:
2153 case ParsedAttr::AT_HLSLIsCounter:
2157 case ParsedAttr::AT_HLSLContainedType: {
2159 Diag(AL.
getLoc(), diag::err_attribute_wrong_number_arguments) << AL << 1;
2165 assert(TSI &&
"no type source info for attribute argument");
2167 diag::err_incomplete_type))
2169 A = HLSLContainedTypeAttr::Create(
getASTContext(), TSI, ACI);
2174 llvm_unreachable(
"unhandled HLSL attribute");
2177 HLSLResourcesTypeAttrs.emplace_back(A);
2183 if (!HLSLResourcesTypeAttrs.size())
2189 HLSLResourcesTypeAttrs, QT, &LocInfo)) {
2190 const HLSLAttributedResourceType *RT =
2197 LocsForHLSLAttributedResources.insert(std::pair(RT, LocInfo));
2199 HLSLResourcesTypeAttrs.clear();
2207 auto I = LocsForHLSLAttributedResources.find(RT);
2208 if (I != LocsForHLSLAttributedResources.end()) {
2209 LocInfo = I->second;
2210 LocsForHLSLAttributedResources.erase(I);
2219void SemaHLSL::collectResourceBindingsOnUserRecordDecl(
const VarDecl *VD,
2220 const RecordType *RT) {
2221 const RecordDecl *RD = RT->getDecl()->getDefinitionOrSelf();
2228 "incomplete arrays inside user defined types are not supported");
2237 if (
const HLSLAttributedResourceType *AttrResType =
2238 HLSLAttributedResourceType::findHandleTypeOnResource(Ty)) {
2243 Bindings.addDeclBindingInfo(VD, RC);
2244 }
else if (
const RecordType *RT = dyn_cast<RecordType>(Ty)) {
2250 collectResourceBindingsOnUserRecordDecl(VD, RT);
2262 bool SpecifiedSpace) {
2263 int RegTypeNum =
static_cast<int>(RegType);
2266 if (D->
hasAttr<HLSLGroupSharedAddressSpaceAttr>()) {
2267 S.
Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
2272 if (
HLSLBufferDecl *CBufferOrTBuffer = dyn_cast<HLSLBufferDecl>(D)) {
2273 ResourceClass RC = CBufferOrTBuffer->isCBuffer() ? ResourceClass::CBuffer
2274 : ResourceClass::SRV;
2284 assert(
isa<VarDecl>(D) &&
"D is expected to be VarDecl or HLSLBufferDecl");
2288 if (
const HLSLAttributedResourceType *AttrResType =
2289 HLSLAttributedResourceType::findHandleTypeOnResource(
2306 if (SpecifiedSpace && !DeclaredInCOrTBuffer)
2307 S.
Diag(ArgLoc, diag::err_hlsl_space_on_global_constant);
2312 if (RegType == RegisterType::CBuffer)
2313 S.
Diag(ArgLoc, diag::warn_hlsl_deprecated_register_type_b);
2314 else if (RegType != RegisterType::C)
2315 S.
Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
2319 if (RegType == RegisterType::C)
2320 S.
Diag(ArgLoc, diag::warn_hlsl_register_type_c_packoffset);
2322 S.
Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
2332 S.
Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
2340 bool RegisterTypesDetected[5] = {
false};
2341 RegisterTypesDetected[
static_cast<int>(regType)] =
true;
2344 if (HLSLResourceBindingAttr *
attr =
2345 dyn_cast<HLSLResourceBindingAttr>(*it)) {
2348 if (RegisterTypesDetected[
static_cast<int>(otherRegType)]) {
2349 int otherRegTypeNum =
static_cast<int>(otherRegType);
2351 diag::err_hlsl_duplicate_register_annotation)
2355 RegisterTypesDetected[
static_cast<int>(otherRegType)] =
true;
2363 bool SpecifiedSpace) {
2368 "expecting VarDecl or HLSLBufferDecl");
2380 const uint64_t &Limit,
2383 uint64_t ArrayCount = 1) {
2388 if (StartSlot > Limit)
2392 if (
const auto *AT = dyn_cast<ArrayType>(T)) {
2395 if (
const auto *CAT = dyn_cast<ConstantArrayType>(AT))
2396 Count = CAT->
getSize().getZExtValue();
2400 ArrayCount * Count);
2404 if (
auto ResTy = dyn_cast<HLSLAttributedResourceType>(T)) {
2407 if (ResTy->getAttrs().ResourceClass != ResClass)
2411 uint64_t EndSlot = StartSlot + ArrayCount - 1;
2412 if (EndSlot > Limit)
2416 StartSlot = EndSlot + 1;
2421 if (
const auto *RT = dyn_cast<RecordType>(T)) {
2424 if (
const auto *CXXRD = dyn_cast<CXXRecordDecl>(RD)) {
2427 ResClass, Ctx, ArrayCount))
2434 ResClass, Ctx, ArrayCount))
2448 const uint64_t Limit = UINT32_MAX;
2449 if (SlotNum > Limit)
2454 if (RegTy == RegisterType::C || RegTy == RegisterType::I)
2457 if (
VarDecl *VD = dyn_cast<VarDecl>(TheDecl)) {
2458 uint64_t BaseSlot = SlotNum;
2466 return (BaseSlot > Limit);
2473 return (SlotNum > Limit);
2476 llvm_unreachable(
"unexpected decl type");
2480 if (
VarDecl *VD = dyn_cast<VarDecl>(TheDecl)) {
2482 if (
const auto *IAT = dyn_cast<IncompleteArrayType>(Ty))
2483 Ty = IAT->getElementType();
2485 diag::err_incomplete_type))
2489 StringRef Slot =
"";
2490 StringRef Space =
"";
2494 Diag(AL.
getLoc(), diag::err_attribute_argument_type)
2504 Diag(AL.
getLoc(), diag::err_attribute_argument_type)
2510 SpaceLoc = Loc->
getLoc();
2513 if (Str.starts_with(
"space")) {
2515 SpaceLoc = Loc->
getLoc();
2524 std::optional<unsigned> SlotNum;
2525 unsigned SpaceNum = 0;
2528 if (!Slot.empty()) {
2530 Diag(SlotLoc, diag::err_hlsl_binding_type_invalid) << Slot.substr(0, 1);
2533 if (RegType == RegisterType::I) {
2534 Diag(SlotLoc, diag::warn_hlsl_deprecated_register_type_i);
2537 const StringRef SlotNumStr = Slot.substr(1);
2542 if (SlotNumStr.getAsInteger(10, N)) {
2543 Diag(SlotLoc, diag::err_hlsl_unsupported_register_number);
2551 Diag(SlotLoc, diag::err_hlsl_register_number_too_large);
2560 if (!Space.starts_with(
"space")) {
2561 Diag(SpaceLoc, diag::err_hlsl_expected_space) << Space;
2564 StringRef SpaceNumStr = Space.substr(5);
2565 if (SpaceNumStr.getAsInteger(10, SpaceNum)) {
2566 Diag(SpaceLoc, diag::err_hlsl_expected_space) << Space;
2571 if (SlotNum.has_value())
2576 HLSLResourceBindingAttr *NewAttr =
2577 HLSLResourceBindingAttr::Create(
getASTContext(), Slot, Space, AL);
2579 NewAttr->setBinding(RegType, SlotNum, SpaceNum);
2634 llvm::DenseMap<const FunctionDecl *, unsigned> ScannedDecls;
2638 llvm::Triple::EnvironmentType CurrentShaderEnvironment;
2639 unsigned CurrentShaderStageBit;
2644 bool ReportOnlyShaderStageIssues;
2647 void SetShaderStageContext(llvm::Triple::EnvironmentType ShaderType) {
2648 static_assert(
sizeof(
unsigned) >= 4);
2649 assert(HLSLShaderAttr::isValidShaderType(ShaderType));
2650 assert((
unsigned)(ShaderType - llvm::Triple::Pixel) < 31 &&
2651 "ShaderType is too big for this bitmap");
2654 unsigned bitmapIndex = ShaderType - llvm::Triple::Pixel;
2655 CurrentShaderEnvironment = ShaderType;
2656 CurrentShaderStageBit = (1 << bitmapIndex);
2659 void SetUnknownShaderStageContext() {
2660 CurrentShaderEnvironment = llvm::Triple::UnknownEnvironment;
2661 CurrentShaderStageBit = (1 << 31);
2664 llvm::Triple::EnvironmentType GetCurrentShaderEnvironment()
const {
2665 return CurrentShaderEnvironment;
2668 bool InUnknownShaderStageContext()
const {
2669 return CurrentShaderEnvironment == llvm::Triple::UnknownEnvironment;
2673 void AddToScannedFunctions(
const FunctionDecl *FD) {
2674 unsigned &ScannedStages = ScannedDecls[FD];
2675 ScannedStages |= CurrentShaderStageBit;
2678 unsigned GetScannedStages(
const FunctionDecl *FD) {
return ScannedDecls[FD]; }
2680 bool WasAlreadyScannedInCurrentStage(
const FunctionDecl *FD) {
2681 return WasAlreadyScannedInCurrentStage(GetScannedStages(FD));
2684 bool WasAlreadyScannedInCurrentStage(
unsigned ScannerStages) {
2685 return ScannerStages & CurrentShaderStageBit;
2688 static bool NeverBeenScanned(
unsigned ScannedStages) {
2689 return ScannedStages == 0;
2693 void HandleFunctionOrMethodRef(FunctionDecl *FD, Expr *RefExpr);
2694 void CheckDeclAvailability(NamedDecl *D,
const AvailabilityAttr *AA,
2696 const AvailabilityAttr *FindAvailabilityAttr(
const Decl *D);
2697 bool HasMatchingEnvironmentOrNone(
const AvailabilityAttr *AA);
2700 DiagnoseHLSLAvailability(Sema &SemaRef)
2702 CurrentShaderEnvironment(llvm::Triple::UnknownEnvironment),
2703 CurrentShaderStageBit(0), ReportOnlyShaderStageIssues(
false) {}
2706 void RunOnTranslationUnit(
const TranslationUnitDecl *TU);
2707 void RunOnFunction(
const FunctionDecl *FD);
2709 bool VisitDeclRefExpr(DeclRefExpr *DRE)
override {
2710 FunctionDecl *FD = llvm::dyn_cast<FunctionDecl>(DRE->
getDecl());
2712 HandleFunctionOrMethodRef(FD, DRE);
2716 bool VisitMemberExpr(MemberExpr *ME)
override {
2717 FunctionDecl *FD = llvm::dyn_cast<FunctionDecl>(ME->
getMemberDecl());
2719 HandleFunctionOrMethodRef(FD, ME);
2724void DiagnoseHLSLAvailability::HandleFunctionOrMethodRef(
FunctionDecl *FD,
2727 "expected DeclRefExpr or MemberExpr");
2731 if (FD->
hasBody(FDWithBody)) {
2732 if (!WasAlreadyScannedInCurrentStage(FDWithBody))
2733 DeclsToScan.push_back(FDWithBody);
2738 const AvailabilityAttr *AA = FindAvailabilityAttr(FD);
2740 CheckDeclAvailability(
2744void DiagnoseHLSLAvailability::RunOnTranslationUnit(
2753 DeclContextsToScan.push_back(TU);
2755 while (!DeclContextsToScan.empty()) {
2756 const DeclContext *DC = DeclContextsToScan.pop_back_val();
2757 for (
auto &D : DC->
decls()) {
2764 if (llvm::dyn_cast<NamespaceDecl>(D) || llvm::dyn_cast<ExportDecl>(D)) {
2765 DeclContextsToScan.push_back(llvm::dyn_cast<DeclContext>(D));
2770 const FunctionDecl *FD = llvm::dyn_cast<FunctionDecl>(D);
2775 if (HLSLShaderAttr *ShaderAttr = FD->
getAttr<HLSLShaderAttr>()) {
2776 SetShaderStageContext(ShaderAttr->getType());
2785 for (
const auto *Redecl : FD->
redecls()) {
2786 if (Redecl->isInExportDeclContext()) {
2793 SetUnknownShaderStageContext();
2801void DiagnoseHLSLAvailability::RunOnFunction(
const FunctionDecl *FD) {
2802 assert(DeclsToScan.empty() &&
"DeclsToScan should be empty");
2803 DeclsToScan.push_back(FD);
2805 while (!DeclsToScan.empty()) {
2813 const unsigned ScannedStages = GetScannedStages(FD);
2814 if (WasAlreadyScannedInCurrentStage(ScannedStages))
2817 ReportOnlyShaderStageIssues = !NeverBeenScanned(ScannedStages);
2819 AddToScannedFunctions(FD);
2824bool DiagnoseHLSLAvailability::HasMatchingEnvironmentOrNone(
2825 const AvailabilityAttr *AA) {
2830 llvm::Triple::EnvironmentType CurrentEnv = GetCurrentShaderEnvironment();
2831 if (CurrentEnv == llvm::Triple::UnknownEnvironment)
2834 llvm::Triple::EnvironmentType AttrEnv =
2835 AvailabilityAttr::getEnvironmentType(IIEnvironment->
getName());
2837 return CurrentEnv == AttrEnv;
2840const AvailabilityAttr *
2841DiagnoseHLSLAvailability::FindAvailabilityAttr(
const Decl *D) {
2842 AvailabilityAttr
const *PartialMatch =
nullptr;
2846 for (
const auto *A : D->
attrs()) {
2847 if (
const auto *Avail = dyn_cast<AvailabilityAttr>(A)) {
2848 StringRef AttrPlatform = Avail->getPlatform()->getName();
2849 StringRef TargetPlatform =
2853 if (AttrPlatform == TargetPlatform) {
2855 if (HasMatchingEnvironmentOrNone(Avail))
2857 PartialMatch = Avail;
2861 return PartialMatch;
2866void DiagnoseHLSLAvailability::CheckDeclAvailability(
NamedDecl *D,
2867 const AvailabilityAttr *AA,
2886 if (ReportOnlyShaderStageIssues)
2892 if (InUnknownShaderStageContext())
2897 bool EnvironmentMatches = HasMatchingEnvironmentOrNone(AA);
2898 VersionTuple Introduced = AA->getIntroduced();
2907 llvm::StringRef PlatformName(
2910 llvm::StringRef CurrentEnvStr =
2911 llvm::Triple::getEnvironmentTypeName(GetCurrentShaderEnvironment());
2913 llvm::StringRef AttrEnvStr =
2914 AA->getEnvironment() ? AA->getEnvironment()->getName() :
"";
2915 bool UseEnvironment = !AttrEnvStr.empty();
2917 if (EnvironmentMatches) {
2918 SemaRef.
Diag(
Range.getBegin(), diag::warn_hlsl_availability)
2919 <<
Range << D << PlatformName << Introduced.getAsString()
2920 << UseEnvironment << CurrentEnvStr;
2922 SemaRef.
Diag(
Range.getBegin(), diag::warn_hlsl_availability_unavailable)
2926 SemaRef.
Diag(D->
getLocation(), diag::note_partial_availability_specified_here)
2927 << D << PlatformName << Introduced.getAsString()
2929 << UseEnvironment << AttrEnvStr << CurrentEnvStr;
2936 if (!DefaultCBufferDecls.empty()) {
2939 DefaultCBufferDecls);
2941 getNextImplicitBindingOrderID());
2942 SemaRef.getCurLexicalContext()->addDecl(DefaultCBuffer);
2946 for (
const Decl *VD : DefaultCBufferDecls) {
2947 const HLSLResourceBindingAttr *RBA =
2948 VD->
getAttr<HLSLResourceBindingAttr>();
2949 if (RBA && RBA->hasRegisterSlot() &&
2950 RBA->getRegisterType() == HLSLResourceBindingAttr::RegisterType::C) {
2957 SemaRef.Consumer.HandleTopLevelDecl(DG);
2959 diagnoseAvailabilityViolations(TU);
2969 TI.
getTriple().getEnvironment() != llvm::Triple::EnvironmentType::Library)
2972 DiagnoseHLSLAvailability(
SemaRef).RunOnTranslationUnit(TU);
2979 for (
unsigned I = 1, N = TheCall->
getNumArgs(); I < N; ++I) {
2982 S->
Diag(TheCall->
getBeginLoc(), diag::err_vec_builtin_incompatible_vector)
3007 for (
unsigned I = 0; I < TheCall->
getNumArgs(); ++I) {
3022 if (!BaseType->isFloat32Type())
3023 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3024 << ArgOrdinal << 5 << 0
3036 if (!BaseType->isHalfType() && !BaseType->isFloat32Type())
3037 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3038 << ArgOrdinal << 5 << 0
3044 unsigned ArgIndex) {
3045 auto *Arg = TheCall->
getArg(ArgIndex);
3047 if (Arg->IgnoreCasts()->isModifiableLvalue(S->
Context, &OrigLoc) ==
3050 S->
Diag(OrigLoc, diag::error_hlsl_inout_lvalue) << Arg << 0;
3060 if (VecTy->getElementType()->isDoubleType())
3061 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3062 << ArgOrdinal << 1 << 0 << 1
3072 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3073 << ArgOrdinal << 5 << 1
3082 if (VecTy->getElementType()->isUnsignedIntegerType())
3085 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3086 << ArgOrdinal << 4 << 3 << 0
3095 return S->
Diag(Loc, diag::err_builtin_invalid_arg_type)
3096 << ArgOrdinal << 5 << 3
3102 unsigned ArgOrdinal,
unsigned Width) {
3105 ArgTy = VTy->getElementType();
3107 uint64_t ElementBitCount =
3109 if (ElementBitCount != Width) {
3111 diag::err_integer_incorrect_bit_count)
3112 << Width << ElementBitCount;
3129 unsigned ArgIndex) {
3138 diag::err_typecheck_expect_scalar_or_vector)
3139 << ArgType << Scalar;
3146 QualType Scalar,
unsigned ArgIndex) {
3157 if (
const auto *VTy = ArgType->getAs<
VectorType>()) {
3170 diag::err_typecheck_expect_scalar_or_vector_or_matrix)
3171 << ArgType << Scalar;
3176 unsigned ArgIndex) {
3181 if (!(ArgType->isScalarType() ||
3182 (VTy && VTy->getElementType()->isScalarType()))) {
3184 diag::err_typecheck_expect_any_scalar_or_vector)
3194 unsigned ArgIndex) {
3196 assert(ArgIndex < TheCall->getNumArgs());
3204 diag::err_typecheck_expect_any_scalar_or_vector)
3229 diag::err_typecheck_call_different_arg_types)
3248 Arg1ScalarTy = VTy->getElementType();
3252 Arg2ScalarTy = VTy->getElementType();
3255 S->
Diag(Arg1->
getBeginLoc(), diag::err_hlsl_builtin_scalar_vector_mismatch)
3256 << 1 << TheCall->
getCallee() << Arg1Ty << Arg2Ty;
3266 if (Arg1Length > 0 && Arg0Length != Arg1Length) {
3268 diag::err_typecheck_vector_lengths_not_equal)
3274 if (Arg2Length > 0 && Arg0Length != Arg2Length) {
3276 diag::err_typecheck_vector_lengths_not_equal)
3289 llvm::function_ref<
bool(
const HLSLAttributedResourceType *ResType)> Check =
3293 const HLSLAttributedResourceType *ResTy =
3297 diag::err_typecheck_expect_hlsl_resource)
3301 if (Check && Check(ResTy)) {
3303 diag::err_invalid_hlsl_resource_type)
3311 QualType BaseType,
unsigned ExpectedCount,
3313 unsigned PassedCount = 1;
3315 PassedCount = VecTy->getNumElements();
3317 if (PassedCount != ExpectedCount) {
3320 S->
Diag(Loc, diag::err_typecheck_convert_incompatible)
3332 [](
const HLSLAttributedResourceType *ResType) {
3333 return ResType->getAttrs().ResourceDimension ==
3334 llvm::dxil::ResourceDimension::Unknown;
3340 [](
const HLSLAttributedResourceType *ResType) {
3341 return ResType->getAttrs().ResourceClass !=
3342 llvm::hlsl::ResourceClass::Sampler;
3350 unsigned ExpectedDim =
3367 unsigned NextIdx = 3;
3373 diag::err_typecheck_convert_incompatible)
3381 Expr *ComponentArg = TheCall->
getArg(NextIdx);
3385 diag::err_typecheck_convert_incompatible)
3392 std::optional<llvm::APSInt> ComponentOpt =
3395 int64_t ComponentVal = ComponentOpt->getSExtValue();
3396 if (ComponentVal != 0) {
3399 assert(ComponentVal >= 0 && ComponentVal <= 3 &&
3400 "The component is not in the expected range.");
3402 diag::err_hlsl_gathercmp_invalid_component)
3412 const HLSLAttributedResourceType *ResourceTy =
3415 unsigned ExpectedDim =
3424 assert(ResourceTy->hasContainedType() &&
3425 "Expecting a contained type for resource with a dimension "
3427 QualType ReturnType = ResourceTy->getContainedType();
3431 S.
Diag(TheCall->
getBeginLoc(), diag::err_hlsl_samplecmp_requires_float);
3437 ReturnType = VecTy->getElementType();
3445 unsigned MinArgs, MaxArgs;
3473 const HLSLAttributedResourceType *ResourceTy =
3475 unsigned ExpectedDim =
3478 unsigned NextIdx = 3;
3487 diag::err_typecheck_convert_incompatible)
3522 diag::err_typecheck_convert_incompatible)
3528 assert(ResourceTy->hasContainedType() &&
3529 "Expecting a contained type for resource with a dimension "
3531 QualType ReturnType = ResourceTy->getContainedType();
3534 S.
Diag(TheCall->
getBeginLoc(), diag::err_hlsl_samplecmp_requires_float);
3547 switch (BuiltinID) {
3548 case Builtin::BI__builtin_hlsl_adduint64: {
3549 if (
SemaRef.checkArgCount(TheCall, 2))
3563 if (NumElementsArg != 2 && NumElementsArg != 4) {
3565 << 1 << 64 << NumElementsArg * 32;
3579 case Builtin::BI__builtin_hlsl_resource_getpointer: {
3580 if (
SemaRef.checkArgCount(TheCall, 2) ||
3583 SemaRef.getASTContext().UnsignedIntTy))
3588 QualType ContainedTy = ResourceTy->getContainedType();
3591 ReturnType =
SemaRef.Context.getPointerType(ReturnType);
3597 case Builtin::BI__builtin_hlsl_resource_getpointer_typed: {
3598 if (
SemaRef.checkArgCount(TheCall, 3) ||
3601 SemaRef.getASTContext().UnsignedIntTy))
3606 "expected pointer type for second argument");
3613 diag::err_invalid_use_of_array_type);
3617 ReturnType =
SemaRef.Context.getPointerType(ReturnType);
3622 case Builtin::BI__builtin_hlsl_resource_load_with_status: {
3623 if (
SemaRef.checkArgCount(TheCall, 3) ||
3626 SemaRef.getASTContext().UnsignedIntTy) ||
3628 SemaRef.getASTContext().UnsignedIntTy) ||
3634 QualType ReturnType = ResourceTy->getContainedType();
3639 case Builtin::BI__builtin_hlsl_resource_load_with_status_typed: {
3640 if (
SemaRef.checkArgCount(TheCall, 4) ||
3643 SemaRef.getASTContext().UnsignedIntTy) ||
3645 SemaRef.getASTContext().UnsignedIntTy) ||
3651 "expected pointer type for second argument");
3658 diag::err_invalid_use_of_array_type);
3664 case Builtin::BI__builtin_hlsl_resource_sample:
3666 case Builtin::BI__builtin_hlsl_resource_sample_bias:
3668 case Builtin::BI__builtin_hlsl_resource_sample_grad:
3670 case Builtin::BI__builtin_hlsl_resource_sample_level:
3672 case Builtin::BI__builtin_hlsl_resource_sample_cmp:
3674 case Builtin::BI__builtin_hlsl_resource_sample_cmp_level_zero:
3676 case Builtin::BI__builtin_hlsl_resource_gather:
3678 case Builtin::BI__builtin_hlsl_resource_gather_cmp:
3680 case Builtin::BI__builtin_hlsl_resource_uninitializedhandle: {
3681 assert(TheCall->
getNumArgs() == 1 &&
"expected 1 arg");
3687 case Builtin::BI__builtin_hlsl_resource_handlefrombinding: {
3688 assert(TheCall->
getNumArgs() == 6 &&
"expected 6 args");
3694 case Builtin::BI__builtin_hlsl_resource_handlefromimplicitbinding: {
3695 assert(TheCall->
getNumArgs() == 6 &&
"expected 6 args");
3701 case Builtin::BI__builtin_hlsl_resource_counterhandlefromimplicitbinding: {
3702 assert(TheCall->
getNumArgs() == 3 &&
"expected 3 args");
3705 auto *MainResType = MainHandleTy->
getAs<HLSLAttributedResourceType>();
3706 auto MainAttrs = MainResType->getAttrs();
3707 assert(!MainAttrs.IsCounter &&
"cannot create a counter from a counter");
3708 MainAttrs.IsCounter =
true;
3710 MainResType->getWrappedType(), MainResType->getContainedType(),
3714 TheCall->
setType(CounterHandleTy);
3717 case Builtin::BI__builtin_hlsl_and:
3718 case Builtin::BI__builtin_hlsl_or: {
3719 if (
SemaRef.checkArgCount(TheCall, 2))
3733 case Builtin::BI__builtin_hlsl_all:
3734 case Builtin::BI__builtin_hlsl_any: {
3735 if (
SemaRef.checkArgCount(TheCall, 1))
3741 case Builtin::BI__builtin_hlsl_asdouble: {
3742 if (
SemaRef.checkArgCount(TheCall, 2))
3746 SemaRef.Context.UnsignedIntTy,
3751 SemaRef.Context.UnsignedIntTy,
3760 case Builtin::BI__builtin_hlsl_elementwise_clamp: {
3761 if (
SemaRef.BuiltinElementwiseTernaryMath(
3767 case Builtin::BI__builtin_hlsl_dot: {
3769 if (
SemaRef.BuiltinVectorToScalarMath(TheCall))
3775 case Builtin::BI__builtin_hlsl_elementwise_firstbithigh:
3776 case Builtin::BI__builtin_hlsl_elementwise_firstbitlow: {
3777 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
3787 EltTy = VecTy->getElementType();
3788 ResTy =
SemaRef.Context.getExtVectorType(ResTy, VecTy->getNumElements());
3801 case Builtin::BI__builtin_hlsl_select: {
3802 if (
SemaRef.checkArgCount(TheCall, 3))
3810 if (VTy && VTy->getElementType()->isBooleanType() &&
3815 case Builtin::BI__builtin_hlsl_elementwise_saturate:
3816 case Builtin::BI__builtin_hlsl_elementwise_rcp: {
3817 if (
SemaRef.checkArgCount(TheCall, 1))
3823 diag::err_builtin_invalid_arg_type)
3826 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
3830 case Builtin::BI__builtin_hlsl_elementwise_degrees:
3831 case Builtin::BI__builtin_hlsl_elementwise_radians:
3832 case Builtin::BI__builtin_hlsl_elementwise_rsqrt:
3833 case Builtin::BI__builtin_hlsl_elementwise_frac:
3834 case Builtin::BI__builtin_hlsl_elementwise_ddx_coarse:
3835 case Builtin::BI__builtin_hlsl_elementwise_ddy_coarse:
3836 case Builtin::BI__builtin_hlsl_elementwise_ddx_fine:
3837 case Builtin::BI__builtin_hlsl_elementwise_ddy_fine: {
3838 if (
SemaRef.checkArgCount(TheCall, 1))
3843 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
3847 case Builtin::BI__builtin_hlsl_elementwise_isinf:
3848 case Builtin::BI__builtin_hlsl_elementwise_isnan: {
3849 if (
SemaRef.checkArgCount(TheCall, 1))
3854 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
3859 case Builtin::BI__builtin_hlsl_lerp: {
3860 if (
SemaRef.checkArgCount(TheCall, 3))
3867 if (
SemaRef.BuiltinElementwiseTernaryMath(TheCall))
3871 case Builtin::BI__builtin_hlsl_mad: {
3872 if (
SemaRef.BuiltinElementwiseTernaryMath(
3878 case Builtin::BI__builtin_hlsl_normalize: {
3879 if (
SemaRef.checkArgCount(TheCall, 1))
3890 case Builtin::BI__builtin_hlsl_elementwise_sign: {
3891 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
3899 case Builtin::BI__builtin_hlsl_step: {
3900 if (
SemaRef.checkArgCount(TheCall, 2))
3912 case Builtin::BI__builtin_hlsl_wave_active_all_equal: {
3913 if (
SemaRef.checkArgCount(TheCall, 1))
3927 unsigned NumElts = VecTy->getNumElements();
3937 case Builtin::BI__builtin_hlsl_wave_active_max:
3938 case Builtin::BI__builtin_hlsl_wave_active_min:
3939 case Builtin::BI__builtin_hlsl_wave_active_sum: {
3940 if (
SemaRef.checkArgCount(TheCall, 1))
3955 case Builtin::BI__builtin_elementwise_bitreverse: {
3963 case Builtin::BI__builtin_hlsl_wave_prefix_count_bits: {
3964 if (
SemaRef.checkArgCount(TheCall, 1))
3969 if (!(
ArgType->isScalarType())) {
3971 diag::err_typecheck_expect_any_scalar_or_vector)
3976 if (!(
ArgType->isBooleanType())) {
3978 diag::err_typecheck_expect_any_scalar_or_vector)
3985 case Builtin::BI__builtin_hlsl_wave_read_lane_at: {
3986 if (
SemaRef.checkArgCount(TheCall, 2))
3994 diag::err_typecheck_convert_incompatible)
3995 << ArgTyIndex <<
SemaRef.Context.UnsignedIntTy << 1 << 0 << 0;
4008 case Builtin::BI__builtin_hlsl_wave_get_lane_index: {
4009 if (
SemaRef.checkArgCount(TheCall, 0))
4013 case Builtin::BI__builtin_hlsl_wave_prefix_sum:
4014 case Builtin::BI__builtin_hlsl_wave_prefix_product: {
4015 if (
SemaRef.checkArgCount(TheCall, 1))
4028 case Builtin::BI__builtin_hlsl_elementwise_splitdouble: {
4029 if (
SemaRef.checkArgCount(TheCall, 3))
4044 case Builtin::BI__builtin_hlsl_elementwise_clip: {
4045 if (
SemaRef.checkArgCount(TheCall, 1))
4052 case Builtin::BI__builtin_elementwise_acos:
4053 case Builtin::BI__builtin_elementwise_asin:
4054 case Builtin::BI__builtin_elementwise_atan:
4055 case Builtin::BI__builtin_elementwise_atan2:
4056 case Builtin::BI__builtin_elementwise_ceil:
4057 case Builtin::BI__builtin_elementwise_cos:
4058 case Builtin::BI__builtin_elementwise_cosh:
4059 case Builtin::BI__builtin_elementwise_exp:
4060 case Builtin::BI__builtin_elementwise_exp2:
4061 case Builtin::BI__builtin_elementwise_exp10:
4062 case Builtin::BI__builtin_elementwise_floor:
4063 case Builtin::BI__builtin_elementwise_fmod:
4064 case Builtin::BI__builtin_elementwise_log:
4065 case Builtin::BI__builtin_elementwise_log2:
4066 case Builtin::BI__builtin_elementwise_log10:
4067 case Builtin::BI__builtin_elementwise_pow:
4068 case Builtin::BI__builtin_elementwise_roundeven:
4069 case Builtin::BI__builtin_elementwise_sin:
4070 case Builtin::BI__builtin_elementwise_sinh:
4071 case Builtin::BI__builtin_elementwise_sqrt:
4072 case Builtin::BI__builtin_elementwise_tan:
4073 case Builtin::BI__builtin_elementwise_tanh:
4074 case Builtin::BI__builtin_elementwise_trunc: {
4080 case Builtin::BI__builtin_hlsl_buffer_update_counter: {
4081 assert(TheCall->
getNumArgs() == 2 &&
"expected 2 args");
4082 auto checkResTy = [](
const HLSLAttributedResourceType *ResTy) ->
bool {
4083 return !(ResTy->getAttrs().ResourceClass == ResourceClass::UAV &&
4084 ResTy->getAttrs().RawBuffer && ResTy->hasContainedType());
4089 std::optional<llvm::APSInt> Offset =
4091 if (!Offset.has_value() ||
std::abs(Offset->getExtValue()) != 1) {
4093 diag::err_hlsl_expect_arg_const_int_one_or_neg_one)
4099 case Builtin::BI__builtin_hlsl_elementwise_f16tof32: {
4100 if (
SemaRef.checkArgCount(TheCall, 1))
4111 ArgTy = VTy->getElementType();
4114 diag::err_builtin_invalid_arg_type)
4123 case Builtin::BI__builtin_hlsl_elementwise_f32tof16: {
4124 if (
SemaRef.checkArgCount(TheCall, 1))
4139 WorkList.push_back(BaseTy);
4140 while (!WorkList.empty()) {
4141 QualType T = WorkList.pop_back_val();
4142 T = T.getCanonicalType().getUnqualifiedType();
4143 if (
const auto *AT = dyn_cast<ConstantArrayType>(T)) {
4151 for (uint64_t Ct = 0; Ct < AT->
getZExtSize(); ++Ct)
4152 llvm::append_range(List, ElementFields);
4157 if (
const auto *VT = dyn_cast<VectorType>(T)) {
4158 List.insert(List.end(), VT->getNumElements(), VT->getElementType());
4161 if (
const auto *MT = dyn_cast<ConstantMatrixType>(T)) {
4162 List.insert(List.end(), MT->getNumElementsFlattened(),
4163 MT->getElementType());
4166 if (
const auto *RD = T->getAsCXXRecordDecl()) {
4167 if (RD->isStandardLayout())
4168 RD = RD->getStandardLayoutBaseWithFields();
4172 if (RD->
isUnion() || !RD->isAggregate()) {
4178 for (
const auto *FD : RD->
fields())
4179 if (!FD->isUnnamedBitField())
4180 FieldTypes.push_back(FD->
getType());
4182 std::reverse(FieldTypes.begin(), FieldTypes.end());
4183 llvm::append_range(WorkList, FieldTypes);
4187 if (!RD->isStandardLayout()) {
4189 for (
const auto &
Base : RD->bases())
4190 FieldTypes.push_back(
Base.getType());
4191 std::reverse(FieldTypes.begin(), FieldTypes.end());
4192 llvm::append_range(WorkList, FieldTypes);
4214 if (
SemaRef.Context.getTypeSize(QT) / 8 > 16)
4220 int ArraySize = VT->getNumElements();
4225 QualType ElTy = VT->getElementType();
4229 if (
SemaRef.Context.getTypeSize(QT) / 8 > 16)
4245 if (
SemaRef.getASTContext().hasSameType(T1, T2))
4254 return llvm::equal(T1Types, T2Types,
4256 return SemaRef.IsLayoutCompatible(LHS, RHS);
4265 bool HadError =
false;
4267 for (
unsigned i = 0, e =
New->getNumParams(); i != e; ++i) {
4275 const auto *NDAttr = NewParam->
getAttr<HLSLParamModifierAttr>();
4276 unsigned NSpellingIdx = (NDAttr ? NDAttr->getSpellingListIndex() : 0);
4277 const auto *ODAttr = OldParam->
getAttr<HLSLParamModifierAttr>();
4278 unsigned OSpellingIdx = (ODAttr ? ODAttr->getSpellingListIndex() : 0);
4280 if (NSpellingIdx != OSpellingIdx) {
4282 diag::err_hlsl_param_qualifier_mismatch)
4283 << NDAttr << NewParam;
4299 if (
SemaRef.getASTContext().hasSameUnqualifiedType(SrcTy, DestTy))
4314 llvm_unreachable(
"HLSL doesn't support pointers.");
4317 llvm_unreachable(
"HLSL doesn't support complex types.");
4319 llvm_unreachable(
"HLSL doesn't support fixed point types.");
4321 llvm_unreachable(
"Should have returned before this");
4331 llvm_unreachable(
"HLSL doesn't support complex types.");
4333 llvm_unreachable(
"HLSL doesn't support fixed point types.");
4338 llvm_unreachable(
"HLSL doesn't support pointers.");
4340 llvm_unreachable(
"Should have returned before this");
4346 llvm_unreachable(
"HLSL doesn't support pointers.");
4349 llvm_unreachable(
"HLSL doesn't support fixed point types.");
4353 llvm_unreachable(
"HLSL doesn't support complex types.");
4356 llvm_unreachable(
"Unhandled scalar cast");
4383 for (
unsigned I = 0, Size = DestTypes.size(); I < Size; ++I) {
4384 if (DestTypes[I]->isUnionType())
4416 if (SrcTypes.size() < DestTypes.size())
4419 unsigned SrcSize = SrcTypes.size();
4420 unsigned DstSize = DestTypes.size();
4422 for (I = 0; I < DstSize && I < SrcSize; I++) {
4423 if (SrcTypes[I]->isUnionType() || DestTypes[I]->isUnionType())
4431 for (; I < SrcSize; I++) {
4432 if (SrcTypes[I]->isUnionType())
4439 assert(Param->hasAttr<HLSLParamModifierAttr>() &&
4440 "We should not get here without a parameter modifier expression");
4441 const auto *
Attr = Param->getAttr<HLSLParamModifierAttr>();
4448 << Arg << (IsInOut ? 1 : 0);
4454 QualType Ty = Param->getType().getNonLValueExprType(Ctx);
4461 << Arg << (IsInOut ? 1 : 0);
4473 SemaRef.PerformCopyInitialization(Entity, Param->getBeginLoc(), ArgOpV);
4479 auto *OpV =
new (Ctx)
4484 Res =
SemaRef.ActOnBinOp(
SemaRef.getCurScope(), Param->getBeginLoc(),
4485 tok::equal, ArgOpV, OpV);
4501 "Pointer and reference types cannot be inout or out parameters");
4502 Ty =
SemaRef.getASTContext().getLValueReferenceType(Ty);
4510 bool IsVKPushConstant = IsVulkan && VD->
hasAttr<HLSLVkPushConstantAttr>();
4515 !VD->
hasAttr<HLSLVkConstantIdAttr>() && !IsVKPushConstant &&
4521 if (
Decl->getType().hasAddressSpace())
4524 if (
Decl->getType()->isDependentType())
4537 llvm::Triple::Vulkan;
4538 if (IsVulkan &&
Decl->
hasAttr<HLSLVkPushConstantAttr>()) {
4539 if (HasDeclaredAPushConstant)
4545 HasDeclaredAPushConstant =
true;
4569 if (
SemaRef.RequireCompleteType(
4572 diag::err_typecheck_decl_incomplete_type)) {
4586 DefaultCBufferDecls.push_back(VD);
4591 collectResourceBindingsOnVarDecl(VD);
4593 if (VD->
hasAttr<HLSLVkConstantIdAttr>())
4605 processExplicitBindingsOnDecl(VD);
4615 uint32_t OrderID = getNextImplicitBindingOrderID();
4633 uint32_t OrderID = getNextImplicitBindingOrderID();
4641 if (VD->
hasAttr<HLSLGroupSharedAddressSpaceAttr>())
4648bool SemaHLSL::initGlobalResourceDecl(
VarDecl *VD) {
4650 "expected resource record type");
4666 const char *CreateMethodName;
4667 if (Binding.isExplicit())
4668 CreateMethodName = HasCounter ?
"__createFromBindingWithImplicitCounter"
4669 :
"__createFromBinding";
4671 CreateMethodName = HasCounter
4672 ?
"__createFromImplicitBindingWithImplicitCounter"
4673 :
"__createFromImplicitBinding";
4684 if (Binding.isExplicit()) {
4688 Args.push_back(RegSlot);
4690 uint32_t OrderID = (Binding.hasImplicitOrderID())
4691 ? Binding.getImplicitOrderID()
4692 : getNextImplicitBindingOrderID();
4696 Args.push_back(OrderId);
4699 IntegerLiteral *Space =
4702 Args.push_back(Space);
4705 AST, llvm::APInt(IntTySize, 1), AST.
IntTy, SourceLocation());
4706 Args.push_back(RangeSize);
4709 AST, llvm::APInt(UIntTySize, 0), AST.
UnsignedIntTy, SourceLocation());
4710 Args.push_back(Index);
4712 StringRef VarName = VD->
getName();
4719 Name,
nullptr,
VK_PRValue, FPOptionsOverride());
4720 Args.push_back(NameCast);
4724 uint32_t CounterOrderID = getNextImplicitBindingOrderID();
4725 IntegerLiteral *CounterId =
4728 Args.push_back(CounterId);
4739 AST, NestedNameSpecifierLoc(), SourceLocation(), CreateMethod,
false,
4744 CK_FunctionToPointerDecay, DRE,
nullptr,
VK_PRValue, FPOptionsOverride());
4746 CallExpr *InitExpr =
4748 SourceLocation(), FPOptionsOverride());
4751 SemaRef.CheckCompleteVariableDeclaration(VD);
4755bool SemaHLSL::initGlobalResourceArrayDecl(
VarDecl *VD) {
4757 "expected array of resource records");
4768 ASTContext &AST =
SemaRef.getASTContext();
4771 CXXMethodDecl *CreateMethod =
nullptr;
4774 ResourceBindingAttrs ResourceAttrs(VD);
4775 if (ResourceAttrs.isExplicit())
4778 lookupMethod(
SemaRef, ResourceDecl,
4779 HasCounter ?
"__createFromBindingWithImplicitCounter"
4780 :
"__createFromBinding",
4784 CreateMethod = lookupMethod(
4786 HasCounter ?
"__createFromImplicitBindingWithImplicitCounter"
4787 :
"__createFromImplicitBinding",
4812 return initGlobalResourceDecl(VD);
4814 return initGlobalResourceArrayDecl(VD);
4819std::optional<const DeclBindingInfo *> SemaHLSL::inferGlobalBinding(
Expr *E) {
4820 if (
auto *Ternary = dyn_cast<ConditionalOperator>(E)) {
4821 auto TrueInfo = inferGlobalBinding(Ternary->getTrueExpr());
4822 auto FalseInfo = inferGlobalBinding(Ternary->getFalseExpr());
4823 if (!TrueInfo || !FalseInfo)
4824 return std::nullopt;
4825 if (*TrueInfo != *FalseInfo)
4826 return std::nullopt;
4830 if (
auto *ASE = dyn_cast<ArraySubscriptExpr>(E))
4839 if (
const auto *AttrResType =
4840 HLSLAttributedResourceType::findHandleTypeOnResource(Ty)) {
4842 return Bindings.getDeclBindingInfo(VD, RC);
4849void SemaHLSL::trackLocalResource(
VarDecl *VD,
Expr *E) {
4850 std::optional<const DeclBindingInfo *> ExprBinding = inferGlobalBinding(E);
4853 diag::warn_hlsl_assigning_local_resource_is_not_unique)
4858 if (*ExprBinding ==
nullptr)
4861 auto PrevBinding = Assigns.find(VD);
4862 if (PrevBinding == Assigns.end()) {
4864 Assigns.insert({VD, *ExprBinding});
4869 if (*ExprBinding != PrevBinding->second) {
4871 diag::warn_hlsl_assigning_local_resource_is_not_unique)
4873 SemaRef.Diag(VD->getLocation(), diag::note_var_declared_here) << VD;
4884 "expected LHS to be a resource record or array of resource records");
4885 if (Opc != BO_Assign)
4890 while (
auto *ASE = dyn_cast<ArraySubscriptExpr>(E))
4898 SemaRef.Diag(Loc, diag::err_hlsl_assign_to_global_resource) << VD;
4903 trackLocalResource(VD, RHSExpr);
4911void SemaHLSL::collectResourceBindingsOnVarDecl(
VarDecl *VD) {
4913 "expected global variable that contains HLSL resource");
4916 if (
const HLSLBufferDecl *CBufferOrTBuffer = dyn_cast<HLSLBufferDecl>(VD)) {
4917 Bindings.addDeclBindingInfo(VD, CBufferOrTBuffer->isCBuffer()
4918 ? ResourceClass::CBuffer
4919 : ResourceClass::SRV);
4932 if (
const HLSLAttributedResourceType *AttrResType =
4933 HLSLAttributedResourceType::findHandleTypeOnResource(Ty)) {
4934 Bindings.addDeclBindingInfo(VD, AttrResType->getAttrs().ResourceClass);
4939 if (
const RecordType *RT = dyn_cast<RecordType>(Ty))
4940 collectResourceBindingsOnUserRecordDecl(VD, RT);
4946void SemaHLSL::processExplicitBindingsOnDecl(
VarDecl *VD) {
4949 bool HasBinding =
false;
4950 for (Attr *A : VD->
attrs()) {
4953 if (
auto PA = VD->
getAttr<HLSLVkPushConstantAttr>())
4954 Diag(PA->getLoc(), diag::err_hlsl_attr_incompatible) << A << PA;
4957 HLSLResourceBindingAttr *RBA = dyn_cast<HLSLResourceBindingAttr>(A);
4958 if (!RBA || !RBA->hasRegisterSlot())
4963 assert(RT != RegisterType::I &&
"invalid or obsolete register type should "
4964 "never have an attribute created");
4966 if (RT == RegisterType::C) {
4967 if (Bindings.hasBindingInfoForDecl(VD))
4969 diag::warn_hlsl_user_defined_type_missing_member)
4970 <<
static_cast<int>(RT);
4978 if (DeclBindingInfo *BI = Bindings.getDeclBindingInfo(VD, RC)) {
4983 diag::warn_hlsl_user_defined_type_missing_member)
4984 <<
static_cast<int>(RT);
4992class InitListTransformer {
4996 QualType *DstIt =
nullptr;
4997 Expr **ArgIt =
nullptr;
5003 bool castInitializer(Expr *E) {
5004 assert(DstIt &&
"This should always be something!");
5005 if (DstIt == DestTypes.end()) {
5007 ArgExprs.push_back(E);
5012 DstIt = DestTypes.begin();
5015 Ctx, *DstIt,
false);
5020 ArgExprs.push_back(
Init);
5025 bool buildInitializerListImpl(Expr *E) {
5027 if (
auto *
Init = dyn_cast<InitListExpr>(E)) {
5028 for (
auto *SubInit :
Init->inits())
5029 if (!buildInitializerListImpl(SubInit))
5038 return castInitializer(E);
5040 if (
auto *VecTy = Ty->
getAs<VectorType>()) {
5045 for (uint64_t I = 0; I <
Size; ++I) {
5047 SizeTy, SourceLocation());
5053 if (!castInitializer(ElExpr.
get()))
5058 if (
auto *MTy = Ty->
getAs<ConstantMatrixType>()) {
5059 unsigned Rows = MTy->getNumRows();
5060 unsigned Cols = MTy->getNumColumns();
5061 QualType ElemTy = MTy->getElementType();
5063 for (
unsigned R = 0; R < Rows; ++R) {
5064 for (
unsigned C = 0;
C < Cols; ++
C) {
5077 if (!castInitializer(ElExpr.
get()))
5085 if (
auto *ArrTy = dyn_cast<ConstantArrayType>(Ty.
getTypePtr())) {
5089 for (uint64_t I = 0; I <
Size; ++I) {
5091 SizeTy, SourceLocation());
5096 if (!buildInitializerListImpl(ElExpr.
get()))
5103 llvm::SmallVector<CXXRecordDecl *> RecordDecls;
5104 RecordDecls.push_back(RD);
5105 while (RecordDecls.back()->getNumBases()) {
5106 CXXRecordDecl *D = RecordDecls.back();
5108 "HLSL doesn't support multiple inheritance");
5109 RecordDecls.push_back(
5112 while (!RecordDecls.empty()) {
5113 CXXRecordDecl *RD = RecordDecls.pop_back_val();
5114 for (
auto *FD : RD->
fields()) {
5115 if (FD->isUnnamedBitField())
5123 if (!buildInitializerListImpl(Res.
get()))
5131 Expr *generateInitListsImpl(QualType Ty) {
5132 assert(ArgIt != ArgExprs.end() &&
"Something is off in iteration!");
5136 llvm::SmallVector<Expr *>
Inits;
5142 if (
auto *ATy = Ty->
getAs<VectorType>()) {
5143 ElTy = ATy->getElementType();
5144 Size = ATy->getNumElements();
5145 }
else if (
auto *CMTy = Ty->
getAs<ConstantMatrixType>()) {
5146 ElTy = CMTy->getElementType();
5147 Size = CMTy->getNumElementsFlattened();
5150 ElTy = VTy->getElementType();
5151 Size = VTy->getZExtSize();
5153 for (uint64_t I = 0; I <
Size; ++I)
5154 Inits.push_back(generateInitListsImpl(ElTy));
5157 llvm::SmallVector<CXXRecordDecl *> RecordDecls;
5158 RecordDecls.push_back(RD);
5159 while (RecordDecls.back()->getNumBases()) {
5160 CXXRecordDecl *D = RecordDecls.back();
5162 "HLSL doesn't support multiple inheritance");
5163 RecordDecls.push_back(
5166 while (!RecordDecls.empty()) {
5167 CXXRecordDecl *RD = RecordDecls.pop_back_val();
5168 for (
auto *FD : RD->
fields())
5169 if (!FD->isUnnamedBitField())
5173 auto *NewInit =
new (Ctx) InitListExpr(Ctx,
Inits.front()->getBeginLoc(),
5175 NewInit->setType(Ty);
5180 llvm::SmallVector<QualType, 16> DestTypes;
5181 llvm::SmallVector<Expr *, 16> ArgExprs;
5182 InitListTransformer(Sema &SemaRef,
const InitializedEntity &Entity)
5183 : S(SemaRef), Ctx(SemaRef.getASTContext()),
5184 Wrap(Entity.
getType()->isIncompleteArrayType()) {
5185 InitTy = Entity.
getType().getNonReferenceType();
5195 DstIt = DestTypes.begin();
5198 bool buildInitializerList(Expr *E) {
return buildInitializerListImpl(E); }
5200 Expr *generateInitLists() {
5201 assert(!ArgExprs.empty() &&
5202 "Call buildInitializerList to generate argument expressions.");
5203 ArgIt = ArgExprs.begin();
5205 return generateInitListsImpl(InitTy);
5206 llvm::SmallVector<Expr *>
Inits;
5207 while (ArgIt != ArgExprs.end())
5208 Inits.push_back(generateInitListsImpl(InitTy));
5210 auto *NewInit =
new (Ctx) InitListExpr(Ctx,
Inits.front()->getBeginLoc(),
5212 llvm::APInt ArySize(64,
Inits.size());
5214 ArraySizeModifier::Normal, 0));
5226 if (
const ArrayType *AT = dyn_cast<ArrayType>(Ty)) {
5233 if (
const auto *RT = Ty->
getAs<RecordType>()) {
5237 if (
const auto *CXXRD = dyn_cast<CXXRecordDecl>(RD)) {
5257 if (
Init->getType()->isScalarType())
5260 InitListTransformer ILT(
SemaRef, Entity);
5262 for (
unsigned I = 0; I <
Init->getNumInits(); ++I) {
5270 Init->setInit(I, E);
5272 if (!ILT.buildInitializerList(E))
5275 size_t ExpectedSize = ILT.DestTypes.size();
5276 size_t ActualSize = ILT.ArgExprs.size();
5277 if (ExpectedSize == 0 && ActualSize == 0)
5284 InitTy =
SemaRef.getASTContext().removeAddrSpaceQualType(InitTy);
5286 SemaRef.Diag(
Init->getBeginLoc(), diag::err_hlsl_incorrect_num_initializers)
5287 << (int)(ExpectedSize < ActualSize) << InitTy
5288 << ExpectedSize << ActualSize;
5298 assert(ExpectedSize > 0 &&
5299 "The expected size of an incomplete array type must be at least 1.");
5301 ((ActualSize + ExpectedSize - 1) / ExpectedSize) * ExpectedSize;
5309 InitTy =
SemaRef.getASTContext().removeAddrSpaceQualType(InitTy);
5310 if (ExpectedSize != ActualSize) {
5311 int TooManyOrFew = ActualSize > ExpectedSize ? 1 : 0;
5312 SemaRef.Diag(
Init->getBeginLoc(), diag::err_hlsl_incorrect_num_initializers)
5313 << TooManyOrFew << InitTy << ExpectedSize << ActualSize;
5320 Init->resizeInits(Ctx, NewInit->getNumInits());
5321 for (
unsigned I = 0; I < NewInit->getNumInits(); ++I)
5322 Init->updateInit(Ctx, I, NewInit->getInit(I));
5330 S.
Diag(OpLoc, diag::err_builtin_matrix_invalid_member)
5340 StringRef AccessorName = CompName->
getName();
5341 assert(!AccessorName.empty() &&
"Matrix Accessor must have a name");
5343 unsigned Rows = MT->getNumRows();
5344 unsigned Cols = MT->getNumColumns();
5345 bool IsZeroBasedAccessor =
false;
5346 unsigned ChunkLen = 0;
5347 if (AccessorName.size() < 2)
5349 "length 4 for zero based: \'_mRC\' or "
5350 "length 3 for one-based: \'_RC\' accessor",
5353 if (AccessorName[0] ==
'_') {
5354 if (AccessorName[1] ==
'm') {
5355 IsZeroBasedAccessor =
true;
5362 S, AccessorName,
"zero based: \'_mRC\' or one-based: \'_RC\' accessor",
5365 if (AccessorName.size() % ChunkLen != 0) {
5366 const llvm::StringRef
Expected = IsZeroBasedAccessor
5367 ?
"zero based: '_mRC' accessor"
5368 :
"one-based: '_RC' accessor";
5373 auto isDigit = [](
char c) {
return c >=
'0' &&
c <=
'9'; };
5374 auto isZeroBasedIndex = [](
unsigned i) {
return i <= 3; };
5375 auto isOneBasedIndex = [](
unsigned i) {
return i >= 1 && i <= 4; };
5377 bool HasRepeated =
false;
5379 unsigned NumComponents = 0;
5380 const char *Begin = AccessorName.data();
5382 for (
unsigned I = 0, E = AccessorName.size(); I < E; I += ChunkLen) {
5383 const char *Chunk = Begin + I;
5384 char RowChar = 0, ColChar = 0;
5385 if (IsZeroBasedAccessor) {
5387 if (Chunk[0] !=
'_' || Chunk[1] !=
'm') {
5388 char Bad = (Chunk[0] !=
'_') ? Chunk[0] : Chunk[1];
5390 S, StringRef(&Bad, 1),
"\'_m\' prefix",
5397 if (Chunk[0] !=
'_')
5399 S, StringRef(&Chunk[0], 1),
"\'_\' prefix",
5406 bool IsDigitsError =
false;
5408 unsigned BadPos = IsZeroBasedAccessor ? 2 : 1;
5412 IsDigitsError =
true;
5416 unsigned BadPos = IsZeroBasedAccessor ? 3 : 2;
5420 IsDigitsError =
true;
5425 unsigned Row = RowChar -
'0';
5426 unsigned Col = ColChar -
'0';
5428 bool HasIndexingError =
false;
5429 if (IsZeroBasedAccessor) {
5431 if (!isZeroBasedIndex(Row)) {
5432 S.
Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
5434 HasIndexingError =
true;
5436 if (!isZeroBasedIndex(Col)) {
5437 S.
Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
5439 HasIndexingError =
true;
5443 if (!isOneBasedIndex(Row)) {
5444 S.
Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
5446 HasIndexingError =
true;
5448 if (!isOneBasedIndex(Col)) {
5449 S.
Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
5451 HasIndexingError =
true;
5458 if (HasIndexingError)
5464 bool HasBoundsError =
false;
5466 Diag(OpLoc, diag::err_hlsl_matrix_index_out_of_bounds)
5468 HasBoundsError =
true;
5471 Diag(OpLoc, diag::err_hlsl_matrix_index_out_of_bounds)
5473 HasBoundsError =
true;
5478 unsigned FlatIndex = Row * Cols + Col;
5479 if (Seen[FlatIndex])
5481 Seen[FlatIndex] =
true;
5484 if (NumComponents == 0 || NumComponents > 4) {
5485 S.
Diag(OpLoc, diag::err_hlsl_matrix_swizzle_invalid_length)
5490 QualType ElemTy = MT->getElementType();
5491 if (NumComponents == 1)
5497 for (Sema::ExtVectorDeclsType::iterator
5501 if ((*I)->getUnderlyingType() == VT)
5512 trackLocalResource(VDecl,
Init);
5514 const HLSLVkConstantIdAttr *ConstIdAttr =
5515 VDecl->
getAttr<HLSLVkConstantIdAttr>();
5522 if (!
Init->isCXX11ConstantExpr(Context, &InitValue)) {
5532 int ConstantID = ConstIdAttr->getId();
5533 llvm::APInt IDVal(Context.getIntWidth(Context.IntTy), ConstantID);
5535 ConstIdAttr->getLocation());
5539 if (
C->getType()->getCanonicalTypeUnqualified() !=
5543 Context.getTrivialTypeSourceInfo(
5544 Init->getType(),
Init->getExprLoc()),
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 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 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 bool CheckAnyScalarOrVector(Sema *S, CallExpr *TheCall, unsigned ArgIndex)
static IdentifierInfo * getHostLayoutStructName(Sema &S, NamedDecl *BaseDecl, bool MustBeUnique)
static void addImplicitBindingAttrToDecl(Sema &S, Decl *D, RegisterType RT, uint32_t ImplicitBindingOrderID)
static void SetElementTypeAsReturnType(Sema *S, CallExpr *TheCall, QualType ReturnType)
static bool isResourceRecordTypeOrArrayOf(VarDecl *VD)
static unsigned calculateLegacyCbufferSize(const ASTContext &Context, QualType T)
static const HLSLAttributedResourceType * getResourceArrayHandleType(VarDecl *VD)
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)
void createHostLayoutStructForBuffer(Sema &S, HLSLBufferDecl *BufDecl)
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 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],...
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.
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.
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.
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
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 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.
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.
NamedDecl * getFoundDecl() const
Fetch the unique decl found by this 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.
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)
void CheckEntryPoint(FunctionDecl *FD)
void emitLogicalOperatorFixIt(Expr *LHS, Expr *RHS, BinaryOperatorKind Opc)
T * createSemanticAttr(const AttributeCommonInfo &ACI, std::optional< unsigned > Location)
void ActOnEndOfTranslationUnit(TranslationUnitDecl *TU)
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
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
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.
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 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.
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
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.
@ 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
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.
void setCounterImplicitOrderID(unsigned Value) const
bool hasCounterImplicitOrderID() const
void setImplicitOrderID(unsigned Value) const
const SourceLocation & getLocation() const
const llvm::hlsl::rootsig::RootElement & getElement() const