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)
3330 unsigned MinArgs, MaxArgs;
3357 [](
const HLSLAttributedResourceType *ResType) {
3358 return ResType->getAttrs().ResourceDimension ==
3359 llvm::dxil::ResourceDimension::Unknown;
3365 [](
const HLSLAttributedResourceType *ResType) {
3366 return ResType->getAttrs().ResourceClass !=
3367 llvm::hlsl::ResourceClass::Sampler;
3375 unsigned ExpectedDim =
3382 unsigned NextIdx = 3;
3391 diag::err_typecheck_convert_incompatible)
3426 diag::err_typecheck_convert_incompatible)
3432 assert(ResourceTy->hasContainedType() &&
3433 "Expecting a contained type for resource with a dimension "
3435 QualType ReturnType = ResourceTy->getContainedType();
3438 S.
Diag(TheCall->
getBeginLoc(), diag::err_hlsl_samplecmp_requires_float);
3451 switch (BuiltinID) {
3452 case Builtin::BI__builtin_hlsl_adduint64: {
3453 if (
SemaRef.checkArgCount(TheCall, 2))
3467 if (NumElementsArg != 2 && NumElementsArg != 4) {
3469 << 1 << 64 << NumElementsArg * 32;
3483 case Builtin::BI__builtin_hlsl_resource_getpointer: {
3484 if (
SemaRef.checkArgCount(TheCall, 2) ||
3487 SemaRef.getASTContext().UnsignedIntTy))
3492 QualType ContainedTy = ResourceTy->getContainedType();
3495 ReturnType =
SemaRef.Context.getPointerType(ReturnType);
3501 case Builtin::BI__builtin_hlsl_resource_getpointer_typed: {
3502 if (
SemaRef.checkArgCount(TheCall, 3) ||
3505 SemaRef.getASTContext().UnsignedIntTy))
3510 "expected pointer type for second argument");
3517 diag::err_invalid_use_of_array_type);
3521 ReturnType =
SemaRef.Context.getPointerType(ReturnType);
3526 case Builtin::BI__builtin_hlsl_resource_load_with_status: {
3527 if (
SemaRef.checkArgCount(TheCall, 3) ||
3530 SemaRef.getASTContext().UnsignedIntTy) ||
3532 SemaRef.getASTContext().UnsignedIntTy) ||
3538 QualType ReturnType = ResourceTy->getContainedType();
3543 case Builtin::BI__builtin_hlsl_resource_load_with_status_typed: {
3544 if (
SemaRef.checkArgCount(TheCall, 4) ||
3547 SemaRef.getASTContext().UnsignedIntTy) ||
3549 SemaRef.getASTContext().UnsignedIntTy) ||
3555 "expected pointer type for second argument");
3562 diag::err_invalid_use_of_array_type);
3568 case Builtin::BI__builtin_hlsl_resource_sample:
3570 case Builtin::BI__builtin_hlsl_resource_sample_bias:
3572 case Builtin::BI__builtin_hlsl_resource_sample_grad:
3574 case Builtin::BI__builtin_hlsl_resource_sample_level:
3576 case Builtin::BI__builtin_hlsl_resource_sample_cmp:
3578 case Builtin::BI__builtin_hlsl_resource_sample_cmp_level_zero:
3580 case Builtin::BI__builtin_hlsl_resource_uninitializedhandle: {
3581 assert(TheCall->
getNumArgs() == 1 &&
"expected 1 arg");
3587 case Builtin::BI__builtin_hlsl_resource_handlefrombinding: {
3588 assert(TheCall->
getNumArgs() == 6 &&
"expected 6 args");
3594 case Builtin::BI__builtin_hlsl_resource_handlefromimplicitbinding: {
3595 assert(TheCall->
getNumArgs() == 6 &&
"expected 6 args");
3601 case Builtin::BI__builtin_hlsl_resource_counterhandlefromimplicitbinding: {
3602 assert(TheCall->
getNumArgs() == 3 &&
"expected 3 args");
3605 auto *MainResType = MainHandleTy->
getAs<HLSLAttributedResourceType>();
3606 auto MainAttrs = MainResType->getAttrs();
3607 assert(!MainAttrs.IsCounter &&
"cannot create a counter from a counter");
3608 MainAttrs.IsCounter =
true;
3610 MainResType->getWrappedType(), MainResType->getContainedType(),
3614 TheCall->
setType(CounterHandleTy);
3617 case Builtin::BI__builtin_hlsl_and:
3618 case Builtin::BI__builtin_hlsl_or: {
3619 if (
SemaRef.checkArgCount(TheCall, 2))
3633 case Builtin::BI__builtin_hlsl_all:
3634 case Builtin::BI__builtin_hlsl_any: {
3635 if (
SemaRef.checkArgCount(TheCall, 1))
3641 case Builtin::BI__builtin_hlsl_asdouble: {
3642 if (
SemaRef.checkArgCount(TheCall, 2))
3646 SemaRef.Context.UnsignedIntTy,
3651 SemaRef.Context.UnsignedIntTy,
3660 case Builtin::BI__builtin_hlsl_elementwise_clamp: {
3661 if (
SemaRef.BuiltinElementwiseTernaryMath(
3667 case Builtin::BI__builtin_hlsl_dot: {
3669 if (
SemaRef.BuiltinVectorToScalarMath(TheCall))
3675 case Builtin::BI__builtin_hlsl_elementwise_firstbithigh:
3676 case Builtin::BI__builtin_hlsl_elementwise_firstbitlow: {
3677 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
3687 EltTy = VecTy->getElementType();
3688 ResTy =
SemaRef.Context.getExtVectorType(ResTy, VecTy->getNumElements());
3701 case Builtin::BI__builtin_hlsl_select: {
3702 if (
SemaRef.checkArgCount(TheCall, 3))
3710 if (VTy && VTy->getElementType()->isBooleanType() &&
3715 case Builtin::BI__builtin_hlsl_elementwise_saturate:
3716 case Builtin::BI__builtin_hlsl_elementwise_rcp: {
3717 if (
SemaRef.checkArgCount(TheCall, 1))
3723 diag::err_builtin_invalid_arg_type)
3726 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
3730 case Builtin::BI__builtin_hlsl_elementwise_degrees:
3731 case Builtin::BI__builtin_hlsl_elementwise_radians:
3732 case Builtin::BI__builtin_hlsl_elementwise_rsqrt:
3733 case Builtin::BI__builtin_hlsl_elementwise_frac:
3734 case Builtin::BI__builtin_hlsl_elementwise_ddx_coarse:
3735 case Builtin::BI__builtin_hlsl_elementwise_ddy_coarse:
3736 case Builtin::BI__builtin_hlsl_elementwise_ddx_fine:
3737 case Builtin::BI__builtin_hlsl_elementwise_ddy_fine: {
3738 if (
SemaRef.checkArgCount(TheCall, 1))
3743 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
3747 case Builtin::BI__builtin_hlsl_elementwise_isinf:
3748 case Builtin::BI__builtin_hlsl_elementwise_isnan: {
3749 if (
SemaRef.checkArgCount(TheCall, 1))
3754 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
3759 case Builtin::BI__builtin_hlsl_lerp: {
3760 if (
SemaRef.checkArgCount(TheCall, 3))
3767 if (
SemaRef.BuiltinElementwiseTernaryMath(TheCall))
3771 case Builtin::BI__builtin_hlsl_mad: {
3772 if (
SemaRef.BuiltinElementwiseTernaryMath(
3778 case Builtin::BI__builtin_hlsl_normalize: {
3779 if (
SemaRef.checkArgCount(TheCall, 1))
3790 case Builtin::BI__builtin_hlsl_elementwise_sign: {
3791 if (
SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
3799 case Builtin::BI__builtin_hlsl_step: {
3800 if (
SemaRef.checkArgCount(TheCall, 2))
3812 case Builtin::BI__builtin_hlsl_wave_active_max:
3813 case Builtin::BI__builtin_hlsl_wave_active_min:
3814 case Builtin::BI__builtin_hlsl_wave_active_sum: {
3815 if (
SemaRef.checkArgCount(TheCall, 1))
3830 case Builtin::BI__builtin_elementwise_bitreverse: {
3838 case Builtin::BI__builtin_hlsl_wave_prefix_count_bits: {
3839 if (
SemaRef.checkArgCount(TheCall, 1))
3844 if (!(
ArgType->isScalarType())) {
3846 diag::err_typecheck_expect_any_scalar_or_vector)
3851 if (!(
ArgType->isBooleanType())) {
3853 diag::err_typecheck_expect_any_scalar_or_vector)
3860 case Builtin::BI__builtin_hlsl_wave_read_lane_at: {
3861 if (
SemaRef.checkArgCount(TheCall, 2))
3869 diag::err_typecheck_convert_incompatible)
3870 << ArgTyIndex <<
SemaRef.Context.UnsignedIntTy << 1 << 0 << 0;
3883 case Builtin::BI__builtin_hlsl_wave_get_lane_index: {
3884 if (
SemaRef.checkArgCount(TheCall, 0))
3888 case Builtin::BI__builtin_hlsl_wave_prefix_sum:
3889 case Builtin::BI__builtin_hlsl_wave_prefix_product: {
3890 if (
SemaRef.checkArgCount(TheCall, 1))
3903 case Builtin::BI__builtin_hlsl_elementwise_splitdouble: {
3904 if (
SemaRef.checkArgCount(TheCall, 3))
3919 case Builtin::BI__builtin_hlsl_elementwise_clip: {
3920 if (
SemaRef.checkArgCount(TheCall, 1))
3927 case Builtin::BI__builtin_elementwise_acos:
3928 case Builtin::BI__builtin_elementwise_asin:
3929 case Builtin::BI__builtin_elementwise_atan:
3930 case Builtin::BI__builtin_elementwise_atan2:
3931 case Builtin::BI__builtin_elementwise_ceil:
3932 case Builtin::BI__builtin_elementwise_cos:
3933 case Builtin::BI__builtin_elementwise_cosh:
3934 case Builtin::BI__builtin_elementwise_exp:
3935 case Builtin::BI__builtin_elementwise_exp2:
3936 case Builtin::BI__builtin_elementwise_exp10:
3937 case Builtin::BI__builtin_elementwise_floor:
3938 case Builtin::BI__builtin_elementwise_fmod:
3939 case Builtin::BI__builtin_elementwise_log:
3940 case Builtin::BI__builtin_elementwise_log2:
3941 case Builtin::BI__builtin_elementwise_log10:
3942 case Builtin::BI__builtin_elementwise_pow:
3943 case Builtin::BI__builtin_elementwise_roundeven:
3944 case Builtin::BI__builtin_elementwise_sin:
3945 case Builtin::BI__builtin_elementwise_sinh:
3946 case Builtin::BI__builtin_elementwise_sqrt:
3947 case Builtin::BI__builtin_elementwise_tan:
3948 case Builtin::BI__builtin_elementwise_tanh:
3949 case Builtin::BI__builtin_elementwise_trunc: {
3955 case Builtin::BI__builtin_hlsl_buffer_update_counter: {
3956 assert(TheCall->
getNumArgs() == 2 &&
"expected 2 args");
3957 auto checkResTy = [](
const HLSLAttributedResourceType *ResTy) ->
bool {
3958 return !(ResTy->getAttrs().ResourceClass == ResourceClass::UAV &&
3959 ResTy->getAttrs().RawBuffer && ResTy->hasContainedType());
3964 std::optional<llvm::APSInt> Offset =
3966 if (!Offset.has_value() ||
std::abs(Offset->getExtValue()) != 1) {
3968 diag::err_hlsl_expect_arg_const_int_one_or_neg_one)
3974 case Builtin::BI__builtin_hlsl_elementwise_f16tof32: {
3975 if (
SemaRef.checkArgCount(TheCall, 1))
3986 ArgTy = VTy->getElementType();
3989 diag::err_builtin_invalid_arg_type)
3998 case Builtin::BI__builtin_hlsl_elementwise_f32tof16: {
3999 if (
SemaRef.checkArgCount(TheCall, 1))
4014 WorkList.push_back(BaseTy);
4015 while (!WorkList.empty()) {
4016 QualType T = WorkList.pop_back_val();
4017 T = T.getCanonicalType().getUnqualifiedType();
4018 if (
const auto *AT = dyn_cast<ConstantArrayType>(T)) {
4026 for (uint64_t Ct = 0; Ct < AT->
getZExtSize(); ++Ct)
4027 llvm::append_range(List, ElementFields);
4032 if (
const auto *VT = dyn_cast<VectorType>(T)) {
4033 List.insert(List.end(), VT->getNumElements(), VT->getElementType());
4036 if (
const auto *MT = dyn_cast<ConstantMatrixType>(T)) {
4037 List.insert(List.end(), MT->getNumElementsFlattened(),
4038 MT->getElementType());
4041 if (
const auto *RD = T->getAsCXXRecordDecl()) {
4042 if (RD->isStandardLayout())
4043 RD = RD->getStandardLayoutBaseWithFields();
4047 if (RD->
isUnion() || !RD->isAggregate()) {
4053 for (
const auto *FD : RD->
fields())
4054 if (!FD->isUnnamedBitField())
4055 FieldTypes.push_back(FD->
getType());
4057 std::reverse(FieldTypes.begin(), FieldTypes.end());
4058 llvm::append_range(WorkList, FieldTypes);
4062 if (!RD->isStandardLayout()) {
4064 for (
const auto &
Base : RD->bases())
4065 FieldTypes.push_back(
Base.getType());
4066 std::reverse(FieldTypes.begin(), FieldTypes.end());
4067 llvm::append_range(WorkList, FieldTypes);
4089 if (
SemaRef.Context.getTypeSize(QT) / 8 > 16)
4095 int ArraySize = VT->getNumElements();
4100 QualType ElTy = VT->getElementType();
4104 if (
SemaRef.Context.getTypeSize(QT) / 8 > 16)
4120 if (
SemaRef.getASTContext().hasSameType(T1, T2))
4129 return llvm::equal(T1Types, T2Types,
4131 return SemaRef.IsLayoutCompatible(LHS, RHS);
4140 bool HadError =
false;
4142 for (
unsigned i = 0, e =
New->getNumParams(); i != e; ++i) {
4150 const auto *NDAttr = NewParam->
getAttr<HLSLParamModifierAttr>();
4151 unsigned NSpellingIdx = (NDAttr ? NDAttr->getSpellingListIndex() : 0);
4152 const auto *ODAttr = OldParam->
getAttr<HLSLParamModifierAttr>();
4153 unsigned OSpellingIdx = (ODAttr ? ODAttr->getSpellingListIndex() : 0);
4155 if (NSpellingIdx != OSpellingIdx) {
4157 diag::err_hlsl_param_qualifier_mismatch)
4158 << NDAttr << NewParam;
4174 if (
SemaRef.getASTContext().hasSameUnqualifiedType(SrcTy, DestTy))
4189 llvm_unreachable(
"HLSL doesn't support pointers.");
4192 llvm_unreachable(
"HLSL doesn't support complex types.");
4194 llvm_unreachable(
"HLSL doesn't support fixed point types.");
4196 llvm_unreachable(
"Should have returned before this");
4206 llvm_unreachable(
"HLSL doesn't support complex types.");
4208 llvm_unreachable(
"HLSL doesn't support fixed point types.");
4213 llvm_unreachable(
"HLSL doesn't support pointers.");
4215 llvm_unreachable(
"Should have returned before this");
4221 llvm_unreachable(
"HLSL doesn't support pointers.");
4224 llvm_unreachable(
"HLSL doesn't support fixed point types.");
4228 llvm_unreachable(
"HLSL doesn't support complex types.");
4231 llvm_unreachable(
"Unhandled scalar cast");
4258 for (
unsigned I = 0, Size = DestTypes.size(); I < Size; ++I) {
4259 if (DestTypes[I]->isUnionType())
4291 if (SrcTypes.size() < DestTypes.size())
4294 unsigned SrcSize = SrcTypes.size();
4295 unsigned DstSize = DestTypes.size();
4297 for (I = 0; I < DstSize && I < SrcSize; I++) {
4298 if (SrcTypes[I]->isUnionType() || DestTypes[I]->isUnionType())
4306 for (; I < SrcSize; I++) {
4307 if (SrcTypes[I]->isUnionType())
4314 assert(Param->hasAttr<HLSLParamModifierAttr>() &&
4315 "We should not get here without a parameter modifier expression");
4316 const auto *
Attr = Param->getAttr<HLSLParamModifierAttr>();
4323 << Arg << (IsInOut ? 1 : 0);
4329 QualType Ty = Param->getType().getNonLValueExprType(Ctx);
4336 << Arg << (IsInOut ? 1 : 0);
4348 SemaRef.PerformCopyInitialization(Entity, Param->getBeginLoc(), ArgOpV);
4354 auto *OpV =
new (Ctx)
4359 Res =
SemaRef.ActOnBinOp(
SemaRef.getCurScope(), Param->getBeginLoc(),
4360 tok::equal, ArgOpV, OpV);
4376 "Pointer and reference types cannot be inout or out parameters");
4377 Ty =
SemaRef.getASTContext().getLValueReferenceType(Ty);
4385 bool IsVKPushConstant = IsVulkan && VD->
hasAttr<HLSLVkPushConstantAttr>();
4390 !VD->
hasAttr<HLSLVkConstantIdAttr>() && !IsVKPushConstant &&
4396 if (
Decl->getType().hasAddressSpace())
4399 if (
Decl->getType()->isDependentType())
4412 llvm::Triple::Vulkan;
4413 if (IsVulkan &&
Decl->
hasAttr<HLSLVkPushConstantAttr>()) {
4414 if (HasDeclaredAPushConstant)
4420 HasDeclaredAPushConstant =
true;
4444 if (
SemaRef.RequireCompleteType(
4447 diag::err_typecheck_decl_incomplete_type)) {
4461 DefaultCBufferDecls.push_back(VD);
4466 collectResourceBindingsOnVarDecl(VD);
4468 if (VD->
hasAttr<HLSLVkConstantIdAttr>())
4480 processExplicitBindingsOnDecl(VD);
4490 uint32_t OrderID = getNextImplicitBindingOrderID();
4508 uint32_t OrderID = getNextImplicitBindingOrderID();
4518bool SemaHLSL::initGlobalResourceDecl(
VarDecl *VD) {
4520 "expected resource record type");
4536 const char *CreateMethodName;
4537 if (Binding.isExplicit())
4538 CreateMethodName = HasCounter ?
"__createFromBindingWithImplicitCounter"
4539 :
"__createFromBinding";
4541 CreateMethodName = HasCounter
4542 ?
"__createFromImplicitBindingWithImplicitCounter"
4543 :
"__createFromImplicitBinding";
4554 if (Binding.isExplicit()) {
4558 Args.push_back(RegSlot);
4560 uint32_t OrderID = (Binding.hasImplicitOrderID())
4561 ? Binding.getImplicitOrderID()
4562 : getNextImplicitBindingOrderID();
4566 Args.push_back(OrderId);
4569 IntegerLiteral *Space =
4572 Args.push_back(Space);
4575 AST, llvm::APInt(IntTySize, 1), AST.
IntTy, SourceLocation());
4576 Args.push_back(RangeSize);
4579 AST, llvm::APInt(UIntTySize, 0), AST.
UnsignedIntTy, SourceLocation());
4580 Args.push_back(Index);
4582 StringRef VarName = VD->
getName();
4589 Name,
nullptr,
VK_PRValue, FPOptionsOverride());
4590 Args.push_back(NameCast);
4594 uint32_t CounterOrderID = getNextImplicitBindingOrderID();
4595 IntegerLiteral *CounterId =
4598 Args.push_back(CounterId);
4609 AST, NestedNameSpecifierLoc(), SourceLocation(), CreateMethod,
false,
4614 CK_FunctionToPointerDecay, DRE,
nullptr,
VK_PRValue, FPOptionsOverride());
4616 CallExpr *InitExpr =
4618 SourceLocation(), FPOptionsOverride());
4621 SemaRef.CheckCompleteVariableDeclaration(VD);
4625bool SemaHLSL::initGlobalResourceArrayDecl(
VarDecl *VD) {
4627 "expected array of resource records");
4638 ASTContext &AST =
SemaRef.getASTContext();
4641 CXXMethodDecl *CreateMethod =
nullptr;
4644 ResourceBindingAttrs ResourceAttrs(VD);
4645 if (ResourceAttrs.isExplicit())
4648 lookupMethod(
SemaRef, ResourceDecl,
4649 HasCounter ?
"__createFromBindingWithImplicitCounter"
4650 :
"__createFromBinding",
4654 CreateMethod = lookupMethod(
4656 HasCounter ?
"__createFromImplicitBindingWithImplicitCounter"
4657 :
"__createFromImplicitBinding",
4682 return initGlobalResourceDecl(VD);
4684 return initGlobalResourceArrayDecl(VD);
4694 "expected LHS to be a resource record or array of resource records");
4695 if (Opc != BO_Assign)
4700 while (
auto *ASE = dyn_cast<ArraySubscriptExpr>(E))
4708 SemaRef.Diag(Loc, diag::err_hlsl_assign_to_global_resource) << VD;
4719void SemaHLSL::collectResourceBindingsOnVarDecl(
VarDecl *VD) {
4721 "expected global variable that contains HLSL resource");
4724 if (
const HLSLBufferDecl *CBufferOrTBuffer = dyn_cast<HLSLBufferDecl>(VD)) {
4725 Bindings.addDeclBindingInfo(VD, CBufferOrTBuffer->isCBuffer()
4726 ? ResourceClass::CBuffer
4727 : ResourceClass::SRV);
4740 if (
const HLSLAttributedResourceType *AttrResType =
4741 HLSLAttributedResourceType::findHandleTypeOnResource(Ty)) {
4742 Bindings.addDeclBindingInfo(VD, AttrResType->getAttrs().ResourceClass);
4747 if (
const RecordType *RT = dyn_cast<RecordType>(Ty))
4748 collectResourceBindingsOnUserRecordDecl(VD, RT);
4754void SemaHLSL::processExplicitBindingsOnDecl(
VarDecl *VD) {
4757 bool HasBinding =
false;
4758 for (Attr *A : VD->
attrs()) {
4761 if (
auto PA = VD->
getAttr<HLSLVkPushConstantAttr>())
4762 Diag(PA->getLoc(), diag::err_hlsl_attr_incompatible) << A << PA;
4765 HLSLResourceBindingAttr *RBA = dyn_cast<HLSLResourceBindingAttr>(A);
4766 if (!RBA || !RBA->hasRegisterSlot())
4771 assert(RT != RegisterType::I &&
"invalid or obsolete register type should "
4772 "never have an attribute created");
4774 if (RT == RegisterType::C) {
4775 if (Bindings.hasBindingInfoForDecl(VD))
4777 diag::warn_hlsl_user_defined_type_missing_member)
4778 <<
static_cast<int>(RT);
4786 if (DeclBindingInfo *BI = Bindings.getDeclBindingInfo(VD, RC)) {
4791 diag::warn_hlsl_user_defined_type_missing_member)
4792 <<
static_cast<int>(RT);
4800class InitListTransformer {
4804 QualType *DstIt =
nullptr;
4805 Expr **ArgIt =
nullptr;
4811 bool castInitializer(Expr *E) {
4812 assert(DstIt &&
"This should always be something!");
4813 if (DstIt == DestTypes.end()) {
4815 ArgExprs.push_back(E);
4820 DstIt = DestTypes.begin();
4823 Ctx, *DstIt,
false);
4828 ArgExprs.push_back(
Init);
4833 bool buildInitializerListImpl(Expr *E) {
4835 if (
auto *
Init = dyn_cast<InitListExpr>(E)) {
4836 for (
auto *SubInit :
Init->inits())
4837 if (!buildInitializerListImpl(SubInit))
4846 return castInitializer(E);
4848 if (
auto *VecTy = Ty->
getAs<VectorType>()) {
4853 for (uint64_t I = 0; I <
Size; ++I) {
4855 SizeTy, SourceLocation());
4861 if (!castInitializer(ElExpr.
get()))
4866 if (
auto *MTy = Ty->
getAs<ConstantMatrixType>()) {
4867 unsigned Rows = MTy->getNumRows();
4868 unsigned Cols = MTy->getNumColumns();
4869 QualType ElemTy = MTy->getElementType();
4871 for (
unsigned R = 0; R < Rows; ++R) {
4872 for (
unsigned C = 0;
C < Cols; ++
C) {
4885 if (!castInitializer(ElExpr.
get()))
4893 if (
auto *ArrTy = dyn_cast<ConstantArrayType>(Ty.
getTypePtr())) {
4897 for (uint64_t I = 0; I <
Size; ++I) {
4899 SizeTy, SourceLocation());
4904 if (!buildInitializerListImpl(ElExpr.
get()))
4911 llvm::SmallVector<CXXRecordDecl *> RecordDecls;
4912 RecordDecls.push_back(RD);
4913 while (RecordDecls.back()->getNumBases()) {
4914 CXXRecordDecl *D = RecordDecls.back();
4916 "HLSL doesn't support multiple inheritance");
4917 RecordDecls.push_back(
4920 while (!RecordDecls.empty()) {
4921 CXXRecordDecl *RD = RecordDecls.pop_back_val();
4922 for (
auto *FD : RD->
fields()) {
4923 if (FD->isUnnamedBitField())
4931 if (!buildInitializerListImpl(Res.
get()))
4939 Expr *generateInitListsImpl(QualType Ty) {
4940 assert(ArgIt != ArgExprs.end() &&
"Something is off in iteration!");
4944 llvm::SmallVector<Expr *>
Inits;
4950 if (
auto *ATy = Ty->
getAs<VectorType>()) {
4951 ElTy = ATy->getElementType();
4952 Size = ATy->getNumElements();
4953 }
else if (
auto *CMTy = Ty->
getAs<ConstantMatrixType>()) {
4954 ElTy = CMTy->getElementType();
4955 Size = CMTy->getNumElementsFlattened();
4958 ElTy = VTy->getElementType();
4959 Size = VTy->getZExtSize();
4961 for (uint64_t I = 0; I <
Size; ++I)
4962 Inits.push_back(generateInitListsImpl(ElTy));
4965 llvm::SmallVector<CXXRecordDecl *> RecordDecls;
4966 RecordDecls.push_back(RD);
4967 while (RecordDecls.back()->getNumBases()) {
4968 CXXRecordDecl *D = RecordDecls.back();
4970 "HLSL doesn't support multiple inheritance");
4971 RecordDecls.push_back(
4974 while (!RecordDecls.empty()) {
4975 CXXRecordDecl *RD = RecordDecls.pop_back_val();
4976 for (
auto *FD : RD->
fields())
4977 if (!FD->isUnnamedBitField())
4981 auto *NewInit =
new (Ctx) InitListExpr(Ctx,
Inits.front()->getBeginLoc(),
4983 NewInit->setType(Ty);
4988 llvm::SmallVector<QualType, 16> DestTypes;
4989 llvm::SmallVector<Expr *, 16> ArgExprs;
4990 InitListTransformer(Sema &SemaRef,
const InitializedEntity &Entity)
4991 : S(SemaRef), Ctx(SemaRef.getASTContext()),
4992 Wrap(Entity.
getType()->isIncompleteArrayType()) {
4993 InitTy = Entity.
getType().getNonReferenceType();
5003 DstIt = DestTypes.begin();
5006 bool buildInitializerList(Expr *E) {
return buildInitializerListImpl(E); }
5008 Expr *generateInitLists() {
5009 assert(!ArgExprs.empty() &&
5010 "Call buildInitializerList to generate argument expressions.");
5011 ArgIt = ArgExprs.begin();
5013 return generateInitListsImpl(InitTy);
5014 llvm::SmallVector<Expr *>
Inits;
5015 while (ArgIt != ArgExprs.end())
5016 Inits.push_back(generateInitListsImpl(InitTy));
5018 auto *NewInit =
new (Ctx) InitListExpr(Ctx,
Inits.front()->getBeginLoc(),
5020 llvm::APInt ArySize(64,
Inits.size());
5022 ArraySizeModifier::Normal, 0));
5034 if (
const ArrayType *AT = dyn_cast<ArrayType>(Ty)) {
5041 if (
const auto *RT = Ty->
getAs<RecordType>()) {
5045 if (
const auto *CXXRD = dyn_cast<CXXRecordDecl>(RD)) {
5065 if (
Init->getType()->isScalarType())
5068 InitListTransformer ILT(
SemaRef, Entity);
5070 for (
unsigned I = 0; I <
Init->getNumInits(); ++I) {
5078 Init->setInit(I, E);
5080 if (!ILT.buildInitializerList(E))
5083 size_t ExpectedSize = ILT.DestTypes.size();
5084 size_t ActualSize = ILT.ArgExprs.size();
5085 if (ExpectedSize == 0 && ActualSize == 0)
5092 InitTy =
SemaRef.getASTContext().removeAddrSpaceQualType(InitTy);
5094 SemaRef.Diag(
Init->getBeginLoc(), diag::err_hlsl_incorrect_num_initializers)
5095 << (int)(ExpectedSize < ActualSize) << InitTy
5096 << ExpectedSize << ActualSize;
5106 assert(ExpectedSize > 0 &&
5107 "The expected size of an incomplete array type must be at least 1.");
5109 ((ActualSize + ExpectedSize - 1) / ExpectedSize) * ExpectedSize;
5117 InitTy =
SemaRef.getASTContext().removeAddrSpaceQualType(InitTy);
5118 if (ExpectedSize != ActualSize) {
5119 int TooManyOrFew = ActualSize > ExpectedSize ? 1 : 0;
5120 SemaRef.Diag(
Init->getBeginLoc(), diag::err_hlsl_incorrect_num_initializers)
5121 << TooManyOrFew << InitTy << ExpectedSize << ActualSize;
5128 Init->resizeInits(Ctx, NewInit->getNumInits());
5129 for (
unsigned I = 0; I < NewInit->getNumInits(); ++I)
5130 Init->updateInit(Ctx, I, NewInit->getInit(I));
5138 S.
Diag(OpLoc, diag::err_builtin_matrix_invalid_member)
5148 StringRef AccessorName = CompName->
getName();
5149 assert(!AccessorName.empty() &&
"Matrix Accessor must have a name");
5151 unsigned Rows = MT->getNumRows();
5152 unsigned Cols = MT->getNumColumns();
5153 bool IsZeroBasedAccessor =
false;
5154 unsigned ChunkLen = 0;
5155 if (AccessorName.size() < 2)
5157 "length 4 for zero based: \'_mRC\' or "
5158 "length 3 for one-based: \'_RC\' accessor",
5161 if (AccessorName[0] ==
'_') {
5162 if (AccessorName[1] ==
'm') {
5163 IsZeroBasedAccessor =
true;
5170 S, AccessorName,
"zero based: \'_mRC\' or one-based: \'_RC\' accessor",
5173 if (AccessorName.size() % ChunkLen != 0) {
5174 const llvm::StringRef
Expected = IsZeroBasedAccessor
5175 ?
"zero based: '_mRC' accessor"
5176 :
"one-based: '_RC' accessor";
5181 auto isDigit = [](
char c) {
return c >=
'0' &&
c <=
'9'; };
5182 auto isZeroBasedIndex = [](
unsigned i) {
return i <= 3; };
5183 auto isOneBasedIndex = [](
unsigned i) {
return i >= 1 && i <= 4; };
5185 bool HasRepeated =
false;
5187 unsigned NumComponents = 0;
5188 const char *Begin = AccessorName.data();
5190 for (
unsigned I = 0, E = AccessorName.size(); I < E; I += ChunkLen) {
5191 const char *Chunk = Begin + I;
5192 char RowChar = 0, ColChar = 0;
5193 if (IsZeroBasedAccessor) {
5195 if (Chunk[0] !=
'_' || Chunk[1] !=
'm') {
5196 char Bad = (Chunk[0] !=
'_') ? Chunk[0] : Chunk[1];
5198 S, StringRef(&Bad, 1),
"\'_m\' prefix",
5205 if (Chunk[0] !=
'_')
5207 S, StringRef(&Chunk[0], 1),
"\'_\' prefix",
5214 bool IsDigitsError =
false;
5216 unsigned BadPos = IsZeroBasedAccessor ? 2 : 1;
5220 IsDigitsError =
true;
5224 unsigned BadPos = IsZeroBasedAccessor ? 3 : 2;
5228 IsDigitsError =
true;
5233 unsigned Row = RowChar -
'0';
5234 unsigned Col = ColChar -
'0';
5236 bool HasIndexingError =
false;
5237 if (IsZeroBasedAccessor) {
5239 if (!isZeroBasedIndex(Row)) {
5240 S.
Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
5242 HasIndexingError =
true;
5244 if (!isZeroBasedIndex(Col)) {
5245 S.
Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
5247 HasIndexingError =
true;
5251 if (!isOneBasedIndex(Row)) {
5252 S.
Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
5254 HasIndexingError =
true;
5256 if (!isOneBasedIndex(Col)) {
5257 S.
Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
5259 HasIndexingError =
true;
5266 if (HasIndexingError)
5272 bool HasBoundsError =
false;
5274 Diag(OpLoc, diag::err_hlsl_matrix_index_out_of_bounds)
5276 HasBoundsError =
true;
5279 Diag(OpLoc, diag::err_hlsl_matrix_index_out_of_bounds)
5281 HasBoundsError =
true;
5286 unsigned FlatIndex = Row * Cols + Col;
5287 if (Seen[FlatIndex])
5289 Seen[FlatIndex] =
true;
5292 if (NumComponents == 0 || NumComponents > 4) {
5293 S.
Diag(OpLoc, diag::err_hlsl_matrix_swizzle_invalid_length)
5298 QualType ElemTy = MT->getElementType();
5299 if (NumComponents == 1)
5305 for (Sema::ExtVectorDeclsType::iterator
5309 if ((*I)->getUnderlyingType() == VT)
5318 const HLSLVkConstantIdAttr *ConstIdAttr =
5319 VDecl->
getAttr<HLSLVkConstantIdAttr>();
5326 if (!
Init->isCXX11ConstantExpr(Context, &InitValue)) {
5336 int ConstantID = ConstIdAttr->getId();
5337 llvm::APInt IDVal(Context.getIntWidth(Context.IntTy), ConstantID);
5339 ConstIdAttr->getLocation());
5343 if (
C->getType()->getCanonicalTypeUnqualified() !=
5347 Context.getTrivialTypeSourceInfo(
5348 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 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 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...
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