clang 23.0.0git
SemaHLSL.cpp
Go to the documentation of this file.
1//===- SemaHLSL.cpp - Semantic Analysis for HLSL constructs ---------------===//
2//
3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4// See https://llvm.org/LICENSE.txt for license information.
5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6//
7//===----------------------------------------------------------------------===//
8// This implements Semantic Analysis for HLSL constructs.
9//===----------------------------------------------------------------------===//
10
11#include "clang/Sema/SemaHLSL.h"
14#include "clang/AST/Attr.h"
15#include "clang/AST/Decl.h"
16#include "clang/AST/DeclBase.h"
17#include "clang/AST/DeclCXX.h"
20#include "clang/AST/Expr.h"
22#include "clang/AST/Type.h"
23#include "clang/AST/TypeBase.h"
24#include "clang/AST/TypeLoc.h"
28#include "clang/Basic/LLVM.h"
33#include "clang/Sema/Lookup.h"
35#include "clang/Sema/Sema.h"
36#include "clang/Sema/Template.h"
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"
50#include <cmath>
51#include <cstddef>
52#include <iterator>
53#include <utility>
54
55using namespace clang;
56using namespace clang::hlsl;
57using RegisterType = HLSLResourceBindingAttr::RegisterType;
58
60 CXXRecordDecl *StructDecl);
61
63 switch (RC) {
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;
72 }
73 llvm_unreachable("unexpected ResourceClass value");
74}
75
76static RegisterType getRegisterType(const HLSLAttributedResourceType *ResTy) {
77 return getRegisterType(ResTy->getAttrs().ResourceClass);
78}
79
80// Converts the first letter of string Slot to RegisterType.
81// Returns false if the letter does not correspond to a valid register type.
82static bool convertToRegisterType(StringRef Slot, RegisterType *RT) {
83 assert(RT != nullptr);
84 switch (Slot[0]) {
85 case 't':
86 case 'T':
87 *RT = RegisterType::SRV;
88 return true;
89 case 'u':
90 case 'U':
91 *RT = RegisterType::UAV;
92 return true;
93 case 'b':
94 case 'B':
95 *RT = RegisterType::CBuffer;
96 return true;
97 case 's':
98 case 'S':
99 *RT = RegisterType::Sampler;
100 return true;
101 case 'c':
102 case 'C':
103 *RT = RegisterType::C;
104 return true;
105 case 'i':
106 case 'I':
107 *RT = RegisterType::I;
108 return true;
109 default:
110 return false;
111 }
112}
113
115 switch (RT) {
116 case RegisterType::SRV:
117 return 't';
118 case RegisterType::UAV:
119 return 'u';
120 case RegisterType::CBuffer:
121 return 'b';
122 case RegisterType::Sampler:
123 return 's';
124 case RegisterType::C:
125 return 'c';
126 case RegisterType::I:
127 return 'i';
128 }
129 llvm_unreachable("unexpected RegisterType value");
130}
131
133 switch (RT) {
134 case RegisterType::SRV:
135 return ResourceClass::SRV;
136 case RegisterType::UAV:
137 return ResourceClass::UAV;
138 case RegisterType::CBuffer:
139 return ResourceClass::CBuffer;
140 case RegisterType::Sampler:
141 return ResourceClass::Sampler;
142 case RegisterType::C:
143 case RegisterType::I:
144 // Deliberately falling through to the unreachable below.
145 break;
146 }
147 llvm_unreachable("unexpected RegisterType value");
148}
149
151 const auto *BT = dyn_cast<BuiltinType>(Type);
152 if (!BT) {
153 if (!Type->isEnumeralType())
154 return Builtin::NotBuiltin;
155 return Builtin::BI__builtin_get_spirv_spec_constant_int;
156 }
157
158 switch (BT->getKind()) {
159 case BuiltinType::Bool:
160 return Builtin::BI__builtin_get_spirv_spec_constant_bool;
161 case BuiltinType::Short:
162 return Builtin::BI__builtin_get_spirv_spec_constant_short;
163 case BuiltinType::Int:
164 return Builtin::BI__builtin_get_spirv_spec_constant_int;
165 case BuiltinType::LongLong:
166 return Builtin::BI__builtin_get_spirv_spec_constant_longlong;
167 case BuiltinType::UShort:
168 return Builtin::BI__builtin_get_spirv_spec_constant_ushort;
169 case BuiltinType::UInt:
170 return Builtin::BI__builtin_get_spirv_spec_constant_uint;
171 case BuiltinType::ULongLong:
172 return Builtin::BI__builtin_get_spirv_spec_constant_ulonglong;
173 case BuiltinType::Half:
174 return Builtin::BI__builtin_get_spirv_spec_constant_half;
175 case BuiltinType::Float:
176 return Builtin::BI__builtin_get_spirv_spec_constant_float;
177 case BuiltinType::Double:
178 return Builtin::BI__builtin_get_spirv_spec_constant_double;
179 default:
180 return Builtin::NotBuiltin;
181 }
182}
183
184static StringRef createRegisterString(ASTContext &AST, RegisterType RegType,
185 unsigned N) {
187 llvm::raw_svector_ostream OS(Buffer);
188 OS << getRegisterTypeChar(RegType);
189 OS << N;
190 return AST.backupStr(OS.str());
191}
192
194 ResourceClass ResClass) {
195 assert(getDeclBindingInfo(VD, ResClass) == nullptr &&
196 "DeclBindingInfo already added");
197 assert(!hasBindingInfoForDecl(VD) || BindingsList.back().Decl == VD);
198 // VarDecl may have multiple entries for different resource classes.
199 // DeclToBindingListIndex stores the index of the first binding we saw
200 // for this decl. If there are any additional ones then that index
201 // shouldn't be updated.
202 DeclToBindingListIndex.try_emplace(VD, BindingsList.size());
203 return &BindingsList.emplace_back(VD, ResClass);
204}
205
207 ResourceClass ResClass) {
208 auto Entry = DeclToBindingListIndex.find(VD);
209 if (Entry != DeclToBindingListIndex.end()) {
210 for (unsigned Index = Entry->getSecond();
211 Index < BindingsList.size() && BindingsList[Index].Decl == VD;
212 ++Index) {
213 if (BindingsList[Index].ResClass == ResClass)
214 return &BindingsList[Index];
215 }
216 }
217 return nullptr;
218}
219
221 return DeclToBindingListIndex.contains(VD);
222}
223
225
226Decl *SemaHLSL::ActOnStartBuffer(Scope *BufferScope, bool CBuffer,
227 SourceLocation KwLoc, IdentifierInfo *Ident,
228 SourceLocation IdentLoc,
229 SourceLocation LBrace) {
230 // For anonymous namespace, take the location of the left brace.
231 DeclContext *LexicalParent = SemaRef.getCurLexicalContext();
233 getASTContext(), LexicalParent, CBuffer, KwLoc, Ident, IdentLoc, LBrace);
234
235 // if CBuffer is false, then it's a TBuffer
236 auto RC = CBuffer ? llvm::hlsl::ResourceClass::CBuffer
237 : llvm::hlsl::ResourceClass::SRV;
238 Result->addAttr(HLSLResourceClassAttr::CreateImplicit(getASTContext(), RC));
239
240 SemaRef.PushOnScopeChains(Result, BufferScope);
241 SemaRef.PushDeclContext(BufferScope, Result);
242
243 return Result;
244}
245
246static unsigned calculateLegacyCbufferFieldAlign(const ASTContext &Context,
247 QualType T) {
248 // Arrays, Matrices, and Structs are always aligned to new buffer rows
249 if (T->isArrayType() || T->isStructureType() || T->isConstantMatrixType())
250 return 16;
251
252 // Vectors are aligned to the type they contain
253 if (const VectorType *VT = T->getAs<VectorType>())
254 return calculateLegacyCbufferFieldAlign(Context, VT->getElementType());
255
256 assert(Context.getTypeSize(T) <= 64 &&
257 "Scalar bit widths larger than 64 not supported");
258
259 // Scalar types are aligned to their byte width
260 return Context.getTypeSize(T) / 8;
261}
262
263// Calculate the size of a legacy cbuffer type in bytes based on
264// https://learn.microsoft.com/en-us/windows/win32/direct3dhlsl/dx-graphics-hlsl-packing-rules
265static unsigned calculateLegacyCbufferSize(const ASTContext &Context,
266 QualType T) {
267 constexpr unsigned CBufferAlign = 16;
268 if (const auto *RD = T->getAsRecordDecl()) {
269 unsigned Size = 0;
270 for (const FieldDecl *Field : RD->fields()) {
271 QualType Ty = Field->getType();
272 unsigned FieldSize = calculateLegacyCbufferSize(Context, Ty);
273 unsigned FieldAlign = calculateLegacyCbufferFieldAlign(Context, Ty);
274
275 // If the field crosses the row boundary after alignment it drops to the
276 // next row
277 unsigned AlignSize = llvm::alignTo(Size, FieldAlign);
278 if ((AlignSize % CBufferAlign) + FieldSize > CBufferAlign) {
279 FieldAlign = CBufferAlign;
280 }
281
282 Size = llvm::alignTo(Size, FieldAlign);
283 Size += FieldSize;
284 }
285 return Size;
286 }
287
288 if (const ConstantArrayType *AT = Context.getAsConstantArrayType(T)) {
289 unsigned ElementCount = AT->getSize().getZExtValue();
290 if (ElementCount == 0)
291 return 0;
292
293 unsigned ElementSize =
294 calculateLegacyCbufferSize(Context, AT->getElementType());
295 unsigned AlignedElementSize = llvm::alignTo(ElementSize, CBufferAlign);
296 return AlignedElementSize * (ElementCount - 1) + ElementSize;
297 }
298
299 if (const VectorType *VT = T->getAs<VectorType>()) {
300 unsigned ElementCount = VT->getNumElements();
301 unsigned ElementSize =
302 calculateLegacyCbufferSize(Context, VT->getElementType());
303 return ElementSize * ElementCount;
304 }
305
306 return Context.getTypeSize(T) / 8;
307}
308
309// Validate packoffset:
310// - if packoffset it used it must be set on all declarations inside the buffer
311// - packoffset ranges must not overlap
312static void validatePackoffset(Sema &S, HLSLBufferDecl *BufDecl) {
314
315 // Make sure the packoffset annotations are either on all declarations
316 // or on none.
317 bool HasPackOffset = false;
318 bool HasNonPackOffset = false;
319 for (auto *Field : BufDecl->buffer_decls()) {
320 VarDecl *Var = dyn_cast<VarDecl>(Field);
321 if (!Var)
322 continue;
323 if (Field->hasAttr<HLSLPackOffsetAttr>()) {
324 PackOffsetVec.emplace_back(Var, Field->getAttr<HLSLPackOffsetAttr>());
325 HasPackOffset = true;
326 } else {
327 HasNonPackOffset = true;
328 }
329 }
330
331 if (!HasPackOffset)
332 return;
333
334 if (HasNonPackOffset)
335 S.Diag(BufDecl->getLocation(), diag::warn_hlsl_packoffset_mix);
336
337 // Make sure there is no overlap in packoffset - sort PackOffsetVec by offset
338 // and compare adjacent values.
339 bool IsValid = true;
340 ASTContext &Context = S.getASTContext();
341 std::sort(PackOffsetVec.begin(), PackOffsetVec.end(),
342 [](const std::pair<VarDecl *, HLSLPackOffsetAttr *> &LHS,
343 const std::pair<VarDecl *, HLSLPackOffsetAttr *> &RHS) {
344 return LHS.second->getOffsetInBytes() <
345 RHS.second->getOffsetInBytes();
346 });
347 for (unsigned i = 0; i < PackOffsetVec.size() - 1; i++) {
348 VarDecl *Var = PackOffsetVec[i].first;
349 HLSLPackOffsetAttr *Attr = PackOffsetVec[i].second;
350 unsigned Size = calculateLegacyCbufferSize(Context, Var->getType());
351 unsigned Begin = Attr->getOffsetInBytes();
352 unsigned End = Begin + Size;
353 unsigned NextBegin = PackOffsetVec[i + 1].second->getOffsetInBytes();
354 if (End > NextBegin) {
355 VarDecl *NextVar = PackOffsetVec[i + 1].first;
356 S.Diag(NextVar->getLocation(), diag::err_hlsl_packoffset_overlap)
357 << NextVar << Var;
358 IsValid = false;
359 }
360 }
361 BufDecl->setHasValidPackoffset(IsValid);
362}
363
364// Returns true if the array has a zero size = if any of the dimensions is 0
365static bool isZeroSizedArray(const ConstantArrayType *CAT) {
366 while (CAT && !CAT->isZeroSize())
367 CAT = dyn_cast<ConstantArrayType>(
369 return CAT != nullptr;
370}
371
375
379
380static const HLSLAttributedResourceType *
382 assert(QT->isHLSLResourceRecordArray() &&
383 "expected array of resource records");
384 const Type *Ty = QT->getUnqualifiedDesugaredType();
385 while (const ArrayType *AT = dyn_cast<ArrayType>(Ty))
387 return HLSLAttributedResourceType::findHandleTypeOnResource(Ty);
388}
389
390static const HLSLAttributedResourceType *
394
395// Returns true if the type is a leaf element type that is not valid to be
396// included in HLSL Buffer, such as a resource class, empty struct, zero-sized
397// array, or a builtin intangible type. Returns false it is a valid leaf element
398// type or if it is a record type that needs to be inspected further.
402 return true;
403 if (const auto *RD = Ty->getAsCXXRecordDecl())
404 return RD->isEmpty();
405 if (Ty->isConstantArrayType() &&
407 return true;
409 return true;
410 return false;
411}
412
413// Returns true if the struct contains at least one element that prevents it
414// from being included inside HLSL Buffer as is, such as an intangible type,
415// empty struct, or zero-sized array. If it does, a new implicit layout struct
416// needs to be created for HLSL Buffer use that will exclude these unwanted
417// declarations (see createHostLayoutStruct function).
419 if (RD->isHLSLIntangible() || RD->isEmpty())
420 return true;
421 // check fields
422 for (const FieldDecl *Field : RD->fields()) {
423 QualType Ty = Field->getType();
425 return true;
426 if (const auto *RD = Ty->getAsCXXRecordDecl();
428 return true;
429 }
430 // check bases
431 for (const CXXBaseSpecifier &Base : RD->bases())
433 Base.getType()->castAsCXXRecordDecl()))
434 return true;
435 return false;
436}
437
439 DeclContext *DC) {
440 CXXRecordDecl *RD = nullptr;
441 for (NamedDecl *Decl :
443 if (CXXRecordDecl *FoundRD = dyn_cast<CXXRecordDecl>(Decl)) {
444 assert(RD == nullptr &&
445 "there should be at most 1 record by a given name in a scope");
446 RD = FoundRD;
447 }
448 }
449 return RD;
450}
451
452// Creates a name for buffer layout struct using the provide name base.
453// If the name must be unique (not previously defined), a suffix is added
454// until a unique name is found.
456 bool MustBeUnique) {
457 ASTContext &AST = S.getASTContext();
458
459 IdentifierInfo *NameBaseII = BaseDecl->getIdentifier();
460 llvm::SmallString<64> Name("__cblayout_");
461 if (NameBaseII) {
462 Name.append(NameBaseII->getName());
463 } else {
464 // anonymous struct
465 Name.append("anon");
466 MustBeUnique = true;
467 }
468
469 size_t NameLength = Name.size();
470 IdentifierInfo *II = &AST.Idents.get(Name, tok::TokenKind::identifier);
471 if (!MustBeUnique)
472 return II;
473
474 unsigned suffix = 0;
475 while (true) {
476 if (suffix != 0) {
477 Name.append("_");
478 Name.append(llvm::Twine(suffix).str());
479 II = &AST.Idents.get(Name, tok::TokenKind::identifier);
480 }
481 if (!findRecordDeclInContext(II, BaseDecl->getDeclContext()))
482 return II;
483 // declaration with that name already exists - increment suffix and try
484 // again until unique name is found
485 suffix++;
486 Name.truncate(NameLength);
487 };
488}
489
490static const Type *createHostLayoutType(Sema &S, const Type *Ty) {
491 ASTContext &AST = S.getASTContext();
492 if (auto *RD = Ty->getAsCXXRecordDecl()) {
494 return Ty;
495 RD = createHostLayoutStruct(S, RD);
496 if (!RD)
497 return nullptr;
498 return AST.getCanonicalTagType(RD)->getTypePtr();
499 }
500
501 if (const auto *CAT = dyn_cast<ConstantArrayType>(Ty)) {
502 const Type *ElementTy = createHostLayoutType(
503 S, CAT->getElementType()->getUnqualifiedDesugaredType());
504 if (!ElementTy)
505 return nullptr;
506 return AST
507 .getConstantArrayType(QualType(ElementTy, 0), CAT->getSize(), nullptr,
508 CAT->getSizeModifier(),
509 CAT->getIndexTypeCVRQualifiers())
510 .getTypePtr();
511 }
512 return Ty;
513}
514
515// Creates a field declaration of given name and type for HLSL buffer layout
516// struct. Returns nullptr if the type cannot be use in HLSL Buffer layout.
518 IdentifierInfo *II,
519 CXXRecordDecl *LayoutStruct) {
521 return nullptr;
522
523 Ty = createHostLayoutType(S, Ty);
524 if (!Ty)
525 return nullptr;
526
527 QualType QT = QualType(Ty, 0);
528 ASTContext &AST = S.getASTContext();
530 auto *Field = FieldDecl::Create(AST, LayoutStruct, SourceLocation(),
531 SourceLocation(), II, QT, TSI, nullptr, false,
533 Field->setAccess(AccessSpecifier::AS_public);
534 return Field;
535}
536
537// Creates host layout struct for a struct included in HLSL Buffer.
538// The layout struct will include only fields that are allowed in HLSL buffer.
539// These fields will be filtered out:
540// - resource classes
541// - empty structs
542// - zero-sized arrays
543// Returns nullptr if the resulting layout struct would be empty.
545 CXXRecordDecl *StructDecl) {
546 assert(requiresImplicitBufferLayoutStructure(StructDecl) &&
547 "struct is already HLSL buffer compatible");
548
549 ASTContext &AST = S.getASTContext();
550 DeclContext *DC = StructDecl->getDeclContext();
551 IdentifierInfo *II = getHostLayoutStructName(S, StructDecl, false);
552
553 // reuse existing if the layout struct if it already exists
554 if (CXXRecordDecl *RD = findRecordDeclInContext(II, DC))
555 return RD;
556
557 CXXRecordDecl *LS =
558 CXXRecordDecl::Create(AST, TagDecl::TagKind::Struct, DC, SourceLocation(),
559 SourceLocation(), II);
560 LS->setImplicit(true);
561 LS->addAttr(PackedAttr::CreateImplicit(AST));
562 LS->startDefinition();
563
564 // copy base struct, create HLSL Buffer compatible version if needed
565 if (unsigned NumBases = StructDecl->getNumBases()) {
566 assert(NumBases == 1 && "HLSL supports only one base type");
567 (void)NumBases;
568 CXXBaseSpecifier Base = *StructDecl->bases_begin();
569 CXXRecordDecl *BaseDecl = Base.getType()->castAsCXXRecordDecl();
571 BaseDecl = createHostLayoutStruct(S, BaseDecl);
572 if (BaseDecl) {
573 TypeSourceInfo *TSI =
575 Base = CXXBaseSpecifier(SourceRange(), false, StructDecl->isClass(),
576 AS_none, TSI, SourceLocation());
577 }
578 }
579 if (BaseDecl) {
580 const CXXBaseSpecifier *BasesArray[1] = {&Base};
581 LS->setBases(BasesArray, 1);
582 }
583 }
584
585 // filter struct fields
586 for (const FieldDecl *FD : StructDecl->fields()) {
587 const Type *Ty = FD->getType()->getUnqualifiedDesugaredType();
588 if (FieldDecl *NewFD =
589 createFieldForHostLayoutStruct(S, Ty, FD->getIdentifier(), LS))
590 LS->addDecl(NewFD);
591 }
592 LS->completeDefinition();
593
594 if (LS->field_empty() && LS->getNumBases() == 0)
595 return nullptr;
596
597 DC->addDecl(LS);
598 return LS;
599}
600
601// Creates host layout struct for HLSL Buffer. The struct will include only
602// fields of types that are allowed in HLSL buffer and it will filter out:
603// - static or groupshared variable declarations
604// - resource classes
605// - empty structs
606// - zero-sized arrays
607// - non-variable declarations
608// The layout struct will be added to the HLSLBufferDecl declarations.
610 ASTContext &AST = S.getASTContext();
611 IdentifierInfo *II = getHostLayoutStructName(S, BufDecl, true);
612
613 CXXRecordDecl *LS =
614 CXXRecordDecl::Create(AST, TagDecl::TagKind::Struct, BufDecl,
616 LS->addAttr(PackedAttr::CreateImplicit(AST));
617 LS->setImplicit(true);
618 LS->startDefinition();
619
620 for (Decl *D : BufDecl->buffer_decls()) {
621 VarDecl *VD = dyn_cast<VarDecl>(D);
622 if (!VD || VD->getStorageClass() == SC_Static ||
624 continue;
625 const Type *Ty = VD->getType()->getUnqualifiedDesugaredType();
626
627 FieldDecl *FD =
629 // Declarations collected for the default $Globals constant buffer have
630 // already been checked to have non-empty cbuffer layout, so
631 // createFieldForHostLayoutStruct should always succeed. These declarations
632 // already have their address space set to hlsl_constant.
633 // For declarations in a named cbuffer block
634 // createFieldForHostLayoutStruct can still return nullptr if the type
635 // is empty (does not have a cbuffer layout).
636 assert((FD || VD->getType().getAddressSpace() != LangAS::hlsl_constant) &&
637 "host layout field for $Globals decl failed to be created");
638 if (FD) {
639 // Add the field decl to the layout struct.
640 LS->addDecl(FD);
642 // Update address space of the original decl to hlsl_constant.
643 QualType NewTy =
645 VD->setType(NewTy);
646 }
647 }
648 }
649 LS->completeDefinition();
650 BufDecl->addLayoutStruct(LS);
651}
652
654 uint32_t ImplicitBindingOrderID) {
655 auto *Attr =
656 HLSLResourceBindingAttr::CreateImplicit(S.getASTContext(), "", "0", {});
657 Attr->setBinding(RT, std::nullopt, 0);
658 Attr->setImplicitBindingOrderID(ImplicitBindingOrderID);
659 D->addAttr(Attr);
660}
661
662// Handle end of cbuffer/tbuffer declaration
664 auto *BufDecl = cast<HLSLBufferDecl>(Dcl);
665 BufDecl->setRBraceLoc(RBrace);
666
667 validatePackoffset(SemaRef, BufDecl);
668
670
671 // Handle implicit binding if needed.
672 ResourceBindingAttrs ResourceAttrs(Dcl);
673 if (!ResourceAttrs.isExplicit()) {
674 SemaRef.Diag(Dcl->getLocation(), diag::warn_hlsl_implicit_binding);
675 // Use HLSLResourceBindingAttr to transfer implicit binding order_ID
676 // to codegen. If it does not exist, create an implicit attribute.
677 uint32_t OrderID = getNextImplicitBindingOrderID();
678 if (ResourceAttrs.hasBinding())
679 ResourceAttrs.setImplicitOrderID(OrderID);
680 else
682 BufDecl->isCBuffer() ? RegisterType::CBuffer
683 : RegisterType::SRV,
684 OrderID);
685 }
686
687 SemaRef.PopDeclContext();
688}
689
690HLSLNumThreadsAttr *SemaHLSL::mergeNumThreadsAttr(Decl *D,
691 const AttributeCommonInfo &AL,
692 int X, int Y, int Z) {
693 if (HLSLNumThreadsAttr *NT = D->getAttr<HLSLNumThreadsAttr>()) {
694 if (NT->getX() != X || NT->getY() != Y || NT->getZ() != Z) {
695 Diag(NT->getLocation(), diag::err_hlsl_attribute_param_mismatch) << AL;
696 Diag(AL.getLoc(), diag::note_conflicting_attribute);
697 }
698 return nullptr;
699 }
700 return ::new (getASTContext())
701 HLSLNumThreadsAttr(getASTContext(), AL, X, Y, Z);
702}
703
705 const AttributeCommonInfo &AL,
706 int Min, int Max, int Preferred,
707 int SpelledArgsCount) {
708 if (HLSLWaveSizeAttr *WS = D->getAttr<HLSLWaveSizeAttr>()) {
709 if (WS->getMin() != Min || WS->getMax() != Max ||
710 WS->getPreferred() != Preferred ||
711 WS->getSpelledArgsCount() != SpelledArgsCount) {
712 Diag(WS->getLocation(), diag::err_hlsl_attribute_param_mismatch) << AL;
713 Diag(AL.getLoc(), diag::note_conflicting_attribute);
714 }
715 return nullptr;
716 }
717 HLSLWaveSizeAttr *Result = ::new (getASTContext())
718 HLSLWaveSizeAttr(getASTContext(), AL, Min, Max, Preferred);
719 Result->setSpelledArgsCount(SpelledArgsCount);
720 return Result;
721}
722
723HLSLVkConstantIdAttr *
725 int Id) {
726
728 if (TargetInfo.getTriple().getArch() != llvm::Triple::spirv) {
729 Diag(AL.getLoc(), diag::warn_attribute_ignored) << AL;
730 return nullptr;
731 }
732
733 auto *VD = cast<VarDecl>(D);
734
735 if (getSpecConstBuiltinId(VD->getType()->getUnqualifiedDesugaredType()) ==
737 Diag(VD->getLocation(), diag::err_specialization_const);
738 return nullptr;
739 }
740
741 if (!VD->getType().isConstQualified()) {
742 Diag(VD->getLocation(), diag::err_specialization_const);
743 return nullptr;
744 }
745
746 if (HLSLVkConstantIdAttr *CI = D->getAttr<HLSLVkConstantIdAttr>()) {
747 if (CI->getId() != Id) {
748 Diag(CI->getLocation(), diag::err_hlsl_attribute_param_mismatch) << AL;
749 Diag(AL.getLoc(), diag::note_conflicting_attribute);
750 }
751 return nullptr;
752 }
753
754 HLSLVkConstantIdAttr *Result =
755 ::new (getASTContext()) HLSLVkConstantIdAttr(getASTContext(), AL, Id);
756 return Result;
757}
758
759HLSLShaderAttr *
761 llvm::Triple::EnvironmentType ShaderType) {
762 if (HLSLShaderAttr *NT = D->getAttr<HLSLShaderAttr>()) {
763 if (NT->getType() != ShaderType) {
764 Diag(NT->getLocation(), diag::err_hlsl_attribute_param_mismatch) << AL;
765 Diag(AL.getLoc(), diag::note_conflicting_attribute);
766 }
767 return nullptr;
768 }
769 return HLSLShaderAttr::Create(getASTContext(), ShaderType, AL);
770}
771
772HLSLParamModifierAttr *
774 HLSLParamModifierAttr::Spelling Spelling) {
775 // We can only merge an `in` attribute with an `out` attribute. All other
776 // combinations of duplicated attributes are ill-formed.
777 if (HLSLParamModifierAttr *PA = D->getAttr<HLSLParamModifierAttr>()) {
778 if ((PA->isIn() && Spelling == HLSLParamModifierAttr::Keyword_out) ||
779 (PA->isOut() && Spelling == HLSLParamModifierAttr::Keyword_in)) {
780 D->dropAttr<HLSLParamModifierAttr>();
781 SourceRange AdjustedRange = {PA->getLocation(), AL.getRange().getEnd()};
782 return HLSLParamModifierAttr::Create(
783 getASTContext(), /*MergedSpelling=*/true, AdjustedRange,
784 HLSLParamModifierAttr::Keyword_inout);
785 }
786 Diag(AL.getLoc(), diag::err_hlsl_duplicate_parameter_modifier) << AL;
787 Diag(PA->getLocation(), diag::note_conflicting_attribute);
788 return nullptr;
789 }
790 return HLSLParamModifierAttr::Create(getASTContext(), AL);
791}
792
795
797 return;
798
799 // If we have specified a root signature to override the entry function then
800 // attach it now
801 HLSLRootSignatureDecl *SignatureDecl =
803 if (SignatureDecl) {
804 FD->dropAttr<RootSignatureAttr>();
805 // We could look up the SourceRange of the macro here as well
806 AttributeCommonInfo AL(RootSigOverrideIdent, AttributeScopeInfo(),
807 SourceRange(), ParsedAttr::Form::Microsoft());
808 FD->addAttr(::new (getASTContext()) RootSignatureAttr(
809 getASTContext(), AL, RootSigOverrideIdent, SignatureDecl));
810 }
811
812 llvm::Triple::EnvironmentType Env = TargetInfo.getTriple().getEnvironment();
813 if (HLSLShaderAttr::isValidShaderType(Env) && Env != llvm::Triple::Library) {
814 if (const auto *Shader = FD->getAttr<HLSLShaderAttr>()) {
815 // The entry point is already annotated - check that it matches the
816 // triple.
817 if (Shader->getType() != Env) {
818 Diag(Shader->getLocation(), diag::err_hlsl_entry_shader_attr_mismatch)
819 << Shader;
820 FD->setInvalidDecl();
821 }
822 } else {
823 // Implicitly add the shader attribute if the entry function isn't
824 // explicitly annotated.
825 FD->addAttr(HLSLShaderAttr::CreateImplicit(getASTContext(), Env,
826 FD->getBeginLoc()));
827 }
828 } else {
829 switch (Env) {
830 case llvm::Triple::UnknownEnvironment:
831 case llvm::Triple::Library:
832 break;
833 case llvm::Triple::RootSignature:
834 llvm_unreachable("rootsig environment has no functions");
835 default:
836 llvm_unreachable("Unhandled environment in triple");
837 }
838 }
839}
840
841static bool isVkPipelineBuiltin(const ASTContext &AstContext, FunctionDecl *FD,
842 HLSLAppliedSemanticAttr *Semantic,
843 bool IsInput) {
844 if (AstContext.getTargetInfo().getTriple().getOS() != llvm::Triple::Vulkan)
845 return false;
846
847 const auto *ShaderAttr = FD->getAttr<HLSLShaderAttr>();
848 assert(ShaderAttr && "Entry point has no shader attribute");
849 llvm::Triple::EnvironmentType ST = ShaderAttr->getType();
850 auto SemanticName = Semantic->getSemanticName().upper();
851
852 // The SV_Position semantic is lowered to:
853 // - Position built-in for vertex output.
854 // - FragCoord built-in for fragment input.
855 if (SemanticName == "SV_POSITION") {
856 return (ST == llvm::Triple::Vertex && !IsInput) ||
857 (ST == llvm::Triple::Pixel && IsInput);
858 }
859 if (SemanticName == "SV_VERTEXID")
860 return true;
861
862 return false;
863}
864
865bool SemaHLSL::determineActiveSemanticOnScalar(FunctionDecl *FD,
866 DeclaratorDecl *OutputDecl,
868 SemanticInfo &ActiveSemantic,
869 SemaHLSL::SemanticContext &SC) {
870 if (ActiveSemantic.Semantic == nullptr) {
871 ActiveSemantic.Semantic = D->getAttr<HLSLParsedSemanticAttr>();
872 if (ActiveSemantic.Semantic)
873 ActiveSemantic.Index = ActiveSemantic.Semantic->getSemanticIndex();
874 }
875
876 if (!ActiveSemantic.Semantic) {
877 Diag(D->getLocation(), diag::err_hlsl_missing_semantic_annotation);
878 return false;
879 }
880
881 auto *A = ::new (getASTContext())
882 HLSLAppliedSemanticAttr(getASTContext(), *ActiveSemantic.Semantic,
883 ActiveSemantic.Semantic->getAttrName()->getName(),
884 ActiveSemantic.Index.value_or(0));
885 if (!A)
887
888 checkSemanticAnnotation(FD, D, A, SC);
889 OutputDecl->addAttr(A);
890
891 unsigned Location = ActiveSemantic.Index.value_or(0);
892
894 SC.CurrentIOType & IOType::In)) {
895 bool HasVkLocation = false;
896 if (auto *A = D->getAttr<HLSLVkLocationAttr>()) {
897 HasVkLocation = true;
898 Location = A->getLocation();
899 }
900
901 if (SC.UsesExplicitVkLocations.value_or(HasVkLocation) != HasVkLocation) {
902 Diag(D->getLocation(), diag::err_hlsl_semantic_partial_explicit_indexing);
903 return false;
904 }
905 SC.UsesExplicitVkLocations = HasVkLocation;
906 }
907
908 const ConstantArrayType *AT = dyn_cast<ConstantArrayType>(D->getType());
909 unsigned ElementCount = AT ? AT->getZExtSize() : 1;
910 ActiveSemantic.Index = Location + ElementCount;
911
912 Twine BaseName = Twine(ActiveSemantic.Semantic->getAttrName()->getName());
913 for (unsigned I = 0; I < ElementCount; ++I) {
914 Twine VariableName = BaseName.concat(Twine(Location + I));
915
916 auto [_, Inserted] = SC.ActiveSemantics.insert(VariableName.str());
917 if (!Inserted) {
918 Diag(D->getLocation(), diag::err_hlsl_semantic_index_overlap)
919 << VariableName.str();
920 return false;
921 }
922 }
923
924 return true;
925}
926
927bool SemaHLSL::determineActiveSemantic(FunctionDecl *FD,
928 DeclaratorDecl *OutputDecl,
930 SemanticInfo &ActiveSemantic,
931 SemaHLSL::SemanticContext &SC) {
932 if (ActiveSemantic.Semantic == nullptr) {
933 ActiveSemantic.Semantic = D->getAttr<HLSLParsedSemanticAttr>();
934 if (ActiveSemantic.Semantic)
935 ActiveSemantic.Index = ActiveSemantic.Semantic->getSemanticIndex();
936 }
937
938 const Type *T = D == FD ? &*FD->getReturnType() : &*D->getType();
940
941 const RecordType *RT = dyn_cast<RecordType>(T);
942 if (!RT)
943 return determineActiveSemanticOnScalar(FD, OutputDecl, D, ActiveSemantic,
944 SC);
945
946 const RecordDecl *RD = RT->getDecl();
947 for (FieldDecl *Field : RD->fields()) {
948 SemanticInfo Info = ActiveSemantic;
949 if (!determineActiveSemantic(FD, OutputDecl, Field, Info, SC)) {
950 Diag(Field->getLocation(), diag::note_hlsl_semantic_used_here) << Field;
951 return false;
952 }
953 if (ActiveSemantic.Semantic)
954 ActiveSemantic = Info;
955 }
956
957 return true;
958}
959
961 const auto *ShaderAttr = FD->getAttr<HLSLShaderAttr>();
962 assert(ShaderAttr && "Entry point has no shader attribute");
963 llvm::Triple::EnvironmentType ST = ShaderAttr->getType();
965 VersionTuple Ver = TargetInfo.getTriple().getOSVersion();
966 switch (ST) {
967 case llvm::Triple::Pixel:
968 case llvm::Triple::Vertex:
969 case llvm::Triple::Geometry:
970 case llvm::Triple::Hull:
971 case llvm::Triple::Domain:
972 case llvm::Triple::RayGeneration:
973 case llvm::Triple::Intersection:
974 case llvm::Triple::AnyHit:
975 case llvm::Triple::ClosestHit:
976 case llvm::Triple::Miss:
977 case llvm::Triple::Callable:
978 if (const auto *NT = FD->getAttr<HLSLNumThreadsAttr>()) {
979 diagnoseAttrStageMismatch(NT, ST,
980 {llvm::Triple::Compute,
981 llvm::Triple::Amplification,
982 llvm::Triple::Mesh});
983 FD->setInvalidDecl();
984 }
985 if (const auto *WS = FD->getAttr<HLSLWaveSizeAttr>()) {
986 diagnoseAttrStageMismatch(WS, ST,
987 {llvm::Triple::Compute,
988 llvm::Triple::Amplification,
989 llvm::Triple::Mesh});
990 FD->setInvalidDecl();
991 }
992 break;
993
994 case llvm::Triple::Compute:
995 case llvm::Triple::Amplification:
996 case llvm::Triple::Mesh:
997 if (!FD->hasAttr<HLSLNumThreadsAttr>()) {
998 Diag(FD->getLocation(), diag::err_hlsl_missing_numthreads)
999 << llvm::Triple::getEnvironmentTypeName(ST);
1000 FD->setInvalidDecl();
1001 }
1002 if (const auto *WS = FD->getAttr<HLSLWaveSizeAttr>()) {
1003 if (Ver < VersionTuple(6, 6)) {
1004 Diag(WS->getLocation(), diag::err_hlsl_attribute_in_wrong_shader_model)
1005 << WS << "6.6";
1006 FD->setInvalidDecl();
1007 } else if (WS->getSpelledArgsCount() > 1 && Ver < VersionTuple(6, 8)) {
1008 Diag(
1009 WS->getLocation(),
1010 diag::err_hlsl_attribute_number_arguments_insufficient_shader_model)
1011 << WS << WS->getSpelledArgsCount() << "6.8";
1012 FD->setInvalidDecl();
1013 }
1014 }
1015 break;
1016 case llvm::Triple::RootSignature:
1017 llvm_unreachable("rootsig environment has no function entry point");
1018 default:
1019 llvm_unreachable("Unhandled environment in triple");
1020 }
1021
1022 SemaHLSL::SemanticContext InputSC = {};
1023 InputSC.CurrentIOType = IOType::In;
1024
1025 for (ParmVarDecl *Param : FD->parameters()) {
1026 SemanticInfo ActiveSemantic;
1027 ActiveSemantic.Semantic = Param->getAttr<HLSLParsedSemanticAttr>();
1028 if (ActiveSemantic.Semantic)
1029 ActiveSemantic.Index = ActiveSemantic.Semantic->getSemanticIndex();
1030
1031 // FIXME: Verify output semantics in parameters.
1032 if (!determineActiveSemantic(FD, Param, Param, ActiveSemantic, InputSC)) {
1033 Diag(Param->getLocation(), diag::note_previous_decl) << Param;
1034 FD->setInvalidDecl();
1035 }
1036 }
1037
1038 SemanticInfo ActiveSemantic;
1039 SemaHLSL::SemanticContext OutputSC = {};
1040 OutputSC.CurrentIOType = IOType::Out;
1041 ActiveSemantic.Semantic = FD->getAttr<HLSLParsedSemanticAttr>();
1042 if (ActiveSemantic.Semantic)
1043 ActiveSemantic.Index = ActiveSemantic.Semantic->getSemanticIndex();
1044 if (!FD->getReturnType()->isVoidType())
1045 determineActiveSemantic(FD, FD, FD, ActiveSemantic, OutputSC);
1046}
1047
1048void SemaHLSL::checkSemanticAnnotation(
1049 FunctionDecl *EntryPoint, const Decl *Param,
1050 const HLSLAppliedSemanticAttr *SemanticAttr, const SemanticContext &SC) {
1051 auto *ShaderAttr = EntryPoint->getAttr<HLSLShaderAttr>();
1052 assert(ShaderAttr && "Entry point has no shader attribute");
1053 llvm::Triple::EnvironmentType ST = ShaderAttr->getType();
1054
1055 auto SemanticName = SemanticAttr->getSemanticName().upper();
1056 if (SemanticName == "SV_DISPATCHTHREADID" ||
1057 SemanticName == "SV_GROUPINDEX" || SemanticName == "SV_GROUPTHREADID" ||
1058 SemanticName == "SV_GROUPID") {
1059
1060 if (ST != llvm::Triple::Compute)
1061 diagnoseSemanticStageMismatch(SemanticAttr, ST, SC.CurrentIOType,
1062 {{llvm::Triple::Compute, IOType::In}});
1063
1064 if (SemanticAttr->getSemanticIndex() != 0) {
1065 std::string PrettyName =
1066 "'" + SemanticAttr->getSemanticName().str() + "'";
1067 Diag(SemanticAttr->getLoc(),
1068 diag::err_hlsl_semantic_indexing_not_supported)
1069 << PrettyName;
1070 }
1071 return;
1072 }
1073
1074 if (SemanticName == "SV_POSITION") {
1075 // SV_Position can be an input or output in vertex shaders,
1076 // but only an input in pixel shaders.
1077 diagnoseSemanticStageMismatch(SemanticAttr, ST, SC.CurrentIOType,
1078 {{llvm::Triple::Vertex, IOType::InOut},
1079 {llvm::Triple::Pixel, IOType::In}});
1080 return;
1081 }
1082 if (SemanticName == "SV_VERTEXID") {
1083 diagnoseSemanticStageMismatch(SemanticAttr, ST, SC.CurrentIOType,
1084 {{llvm::Triple::Vertex, IOType::In}});
1085 return;
1086 }
1087
1088 if (SemanticName == "SV_TARGET") {
1089 diagnoseSemanticStageMismatch(SemanticAttr, ST, SC.CurrentIOType,
1090 {{llvm::Triple::Pixel, IOType::Out}});
1091 return;
1092 }
1093
1094 // FIXME: catch-all for non-implemented system semantics reaching this
1095 // location.
1096 if (SemanticAttr->getAttrName()->getName().starts_with_insensitive("SV_"))
1097 llvm_unreachable("Unknown SemanticAttr");
1098}
1099
1100void SemaHLSL::diagnoseAttrStageMismatch(
1101 const Attr *A, llvm::Triple::EnvironmentType Stage,
1102 std::initializer_list<llvm::Triple::EnvironmentType> AllowedStages) {
1103 SmallVector<StringRef, 8> StageStrings;
1104 llvm::transform(AllowedStages, std::back_inserter(StageStrings),
1105 [](llvm::Triple::EnvironmentType ST) {
1106 return StringRef(
1107 HLSLShaderAttr::ConvertEnvironmentTypeToStr(ST));
1108 });
1109 Diag(A->getLoc(), diag::err_hlsl_attr_unsupported_in_stage)
1110 << A->getAttrName() << llvm::Triple::getEnvironmentTypeName(Stage)
1111 << (AllowedStages.size() != 1) << join(StageStrings, ", ");
1112}
1113
1114void SemaHLSL::diagnoseSemanticStageMismatch(
1115 const Attr *A, llvm::Triple::EnvironmentType Stage, IOType CurrentIOType,
1116 std::initializer_list<SemanticStageInfo> Allowed) {
1117
1118 for (auto &Case : Allowed) {
1119 if (Case.Stage != Stage)
1120 continue;
1121
1122 if (CurrentIOType & Case.AllowedIOTypesMask)
1123 return;
1124
1125 SmallVector<std::string, 8> ValidCases;
1126 llvm::transform(
1127 Allowed, std::back_inserter(ValidCases), [](SemanticStageInfo Case) {
1128 SmallVector<std::string, 2> ValidType;
1129 if (Case.AllowedIOTypesMask & IOType::In)
1130 ValidType.push_back("input");
1131 if (Case.AllowedIOTypesMask & IOType::Out)
1132 ValidType.push_back("output");
1133 return std::string(
1134 HLSLShaderAttr::ConvertEnvironmentTypeToStr(Case.Stage)) +
1135 " " + join(ValidType, "/");
1136 });
1137 Diag(A->getLoc(), diag::err_hlsl_semantic_unsupported_iotype_for_stage)
1138 << A->getAttrName() << (CurrentIOType & IOType::In ? "input" : "output")
1139 << llvm::Triple::getEnvironmentTypeName(Case.Stage)
1140 << join(ValidCases, ", ");
1141 return;
1142 }
1143
1144 SmallVector<StringRef, 8> StageStrings;
1145 llvm::transform(
1146 Allowed, std::back_inserter(StageStrings), [](SemanticStageInfo Case) {
1147 return StringRef(
1148 HLSLShaderAttr::ConvertEnvironmentTypeToStr(Case.Stage));
1149 });
1150
1151 Diag(A->getLoc(), diag::err_hlsl_attr_unsupported_in_stage)
1152 << A->getAttrName() << llvm::Triple::getEnvironmentTypeName(Stage)
1153 << (Allowed.size() != 1) << join(StageStrings, ", ");
1154}
1155
1156template <CastKind Kind>
1157static void castVector(Sema &S, ExprResult &E, QualType &Ty, unsigned Sz) {
1158 if (const auto *VTy = Ty->getAs<VectorType>())
1159 Ty = VTy->getElementType();
1160 Ty = S.getASTContext().getExtVectorType(Ty, Sz);
1161 E = S.ImpCastExprToType(E.get(), Ty, Kind);
1162}
1163
1164template <CastKind Kind>
1166 E = S.ImpCastExprToType(E.get(), Ty, Kind);
1167 return Ty;
1168}
1169
1171 Sema &SemaRef, ExprResult &LHS, ExprResult &RHS, QualType LHSType,
1172 QualType RHSType, QualType LElTy, QualType RElTy, bool IsCompAssign) {
1173 bool LHSFloat = LElTy->isRealFloatingType();
1174 bool RHSFloat = RElTy->isRealFloatingType();
1175
1176 if (LHSFloat && RHSFloat) {
1177 if (IsCompAssign ||
1178 SemaRef.getASTContext().getFloatingTypeOrder(LElTy, RElTy) > 0)
1179 return castElement<CK_FloatingCast>(SemaRef, RHS, LHSType);
1180
1181 return castElement<CK_FloatingCast>(SemaRef, LHS, RHSType);
1182 }
1183
1184 if (LHSFloat)
1185 return castElement<CK_IntegralToFloating>(SemaRef, RHS, LHSType);
1186
1187 assert(RHSFloat);
1188 if (IsCompAssign)
1189 return castElement<clang::CK_FloatingToIntegral>(SemaRef, RHS, LHSType);
1190
1191 return castElement<CK_IntegralToFloating>(SemaRef, LHS, RHSType);
1192}
1193
1195 Sema &SemaRef, ExprResult &LHS, ExprResult &RHS, QualType LHSType,
1196 QualType RHSType, QualType LElTy, QualType RElTy, bool IsCompAssign) {
1197
1198 int IntOrder = SemaRef.Context.getIntegerTypeOrder(LElTy, RElTy);
1199 bool LHSSigned = LElTy->hasSignedIntegerRepresentation();
1200 bool RHSSigned = RElTy->hasSignedIntegerRepresentation();
1201 auto &Ctx = SemaRef.getASTContext();
1202
1203 // If both types have the same signedness, use the higher ranked type.
1204 if (LHSSigned == RHSSigned) {
1205 if (IsCompAssign || IntOrder >= 0)
1206 return castElement<CK_IntegralCast>(SemaRef, RHS, LHSType);
1207
1208 return castElement<CK_IntegralCast>(SemaRef, LHS, RHSType);
1209 }
1210
1211 // If the unsigned type has greater than or equal rank of the signed type, use
1212 // the unsigned type.
1213 if (IntOrder != (LHSSigned ? 1 : -1)) {
1214 if (IsCompAssign || RHSSigned)
1215 return castElement<CK_IntegralCast>(SemaRef, RHS, LHSType);
1216 return castElement<CK_IntegralCast>(SemaRef, LHS, RHSType);
1217 }
1218
1219 // At this point the signed type has higher rank than the unsigned type, which
1220 // means it will be the same size or bigger. If the signed type is bigger, it
1221 // can represent all the values of the unsigned type, so select it.
1222 if (Ctx.getIntWidth(LElTy) != Ctx.getIntWidth(RElTy)) {
1223 if (IsCompAssign || LHSSigned)
1224 return castElement<CK_IntegralCast>(SemaRef, RHS, LHSType);
1225 return castElement<CK_IntegralCast>(SemaRef, LHS, RHSType);
1226 }
1227
1228 // This is a bit of an odd duck case in HLSL. It shouldn't happen, but can due
1229 // to C/C++ leaking through. The place this happens today is long vs long
1230 // long. When arguments are vector<unsigned long, N> and vector<long long, N>,
1231 // the long long has higher rank than long even though they are the same size.
1232
1233 // If this is a compound assignment cast the right hand side to the left hand
1234 // side's type.
1235 if (IsCompAssign)
1236 return castElement<CK_IntegralCast>(SemaRef, RHS, LHSType);
1237
1238 // If this isn't a compound assignment we convert to unsigned long long.
1239 QualType ElTy = Ctx.getCorrespondingUnsignedType(LHSSigned ? LElTy : RElTy);
1240 QualType NewTy = Ctx.getExtVectorType(
1241 ElTy, RHSType->castAs<VectorType>()->getNumElements());
1242 (void)castElement<CK_IntegralCast>(SemaRef, RHS, NewTy);
1243
1244 return castElement<CK_IntegralCast>(SemaRef, LHS, NewTy);
1245}
1246
1248 QualType SrcTy) {
1249 if (DestTy->isRealFloatingType() && SrcTy->isRealFloatingType())
1250 return CK_FloatingCast;
1251 if (DestTy->isIntegralType(Ctx) && SrcTy->isIntegralType(Ctx))
1252 return CK_IntegralCast;
1253 if (DestTy->isRealFloatingType())
1254 return CK_IntegralToFloating;
1255 assert(SrcTy->isRealFloatingType() && DestTy->isIntegralType(Ctx));
1256 return CK_FloatingToIntegral;
1257}
1258
1260 QualType LHSType,
1261 QualType RHSType,
1262 bool IsCompAssign) {
1263 const auto *LVecTy = LHSType->getAs<VectorType>();
1264 const auto *RVecTy = RHSType->getAs<VectorType>();
1265 auto &Ctx = getASTContext();
1266
1267 // If the LHS is not a vector and this is a compound assignment, we truncate
1268 // the argument to a scalar then convert it to the LHS's type.
1269 if (!LVecTy && IsCompAssign) {
1270 QualType RElTy = RHSType->castAs<VectorType>()->getElementType();
1271 RHS = SemaRef.ImpCastExprToType(RHS.get(), RElTy, CK_HLSLVectorTruncation);
1272 RHSType = RHS.get()->getType();
1273 if (Ctx.hasSameUnqualifiedType(LHSType, RHSType))
1274 return LHSType;
1275 RHS = SemaRef.ImpCastExprToType(RHS.get(), LHSType,
1276 getScalarCastKind(Ctx, LHSType, RHSType));
1277 return LHSType;
1278 }
1279
1280 unsigned EndSz = std::numeric_limits<unsigned>::max();
1281 unsigned LSz = 0;
1282 if (LVecTy)
1283 LSz = EndSz = LVecTy->getNumElements();
1284 if (RVecTy)
1285 EndSz = std::min(RVecTy->getNumElements(), EndSz);
1286 assert(EndSz != std::numeric_limits<unsigned>::max() &&
1287 "one of the above should have had a value");
1288
1289 // In a compound assignment, the left operand does not change type, the right
1290 // operand is converted to the type of the left operand.
1291 if (IsCompAssign && LSz != EndSz) {
1292 Diag(LHS.get()->getBeginLoc(),
1293 diag::err_hlsl_vector_compound_assignment_truncation)
1294 << LHSType << RHSType;
1295 return QualType();
1296 }
1297
1298 if (RVecTy && RVecTy->getNumElements() > EndSz)
1299 castVector<CK_HLSLVectorTruncation>(SemaRef, RHS, RHSType, EndSz);
1300 if (!IsCompAssign && LVecTy && LVecTy->getNumElements() > EndSz)
1301 castVector<CK_HLSLVectorTruncation>(SemaRef, LHS, LHSType, EndSz);
1302
1303 if (!RVecTy)
1304 castVector<CK_VectorSplat>(SemaRef, RHS, RHSType, EndSz);
1305 if (!IsCompAssign && !LVecTy)
1306 castVector<CK_VectorSplat>(SemaRef, LHS, LHSType, EndSz);
1307
1308 // If we're at the same type after resizing we can stop here.
1309 if (Ctx.hasSameUnqualifiedType(LHSType, RHSType))
1310 return Ctx.getCommonSugaredType(LHSType, RHSType);
1311
1312 QualType LElTy = LHSType->castAs<VectorType>()->getElementType();
1313 QualType RElTy = RHSType->castAs<VectorType>()->getElementType();
1314
1315 // Handle conversion for floating point vectors.
1316 if (LElTy->isRealFloatingType() || RElTy->isRealFloatingType())
1317 return handleFloatVectorBinOpConversion(SemaRef, LHS, RHS, LHSType, RHSType,
1318 LElTy, RElTy, IsCompAssign);
1319
1320 assert(LElTy->isIntegralType(Ctx) && RElTy->isIntegralType(Ctx) &&
1321 "HLSL Vectors can only contain integer or floating point types");
1322 return handleIntegerVectorBinOpConversion(SemaRef, LHS, RHS, LHSType, RHSType,
1323 LElTy, RElTy, IsCompAssign);
1324}
1325
1327 BinaryOperatorKind Opc) {
1328 assert((Opc == BO_LOr || Opc == BO_LAnd) &&
1329 "Called with non-logical operator");
1331 llvm::raw_svector_ostream OS(Buff);
1332 PrintingPolicy PP(SemaRef.getLangOpts());
1333 StringRef NewFnName = Opc == BO_LOr ? "or" : "and";
1334 OS << NewFnName << "(";
1335 LHS->printPretty(OS, nullptr, PP);
1336 OS << ", ";
1337 RHS->printPretty(OS, nullptr, PP);
1338 OS << ")";
1339 SourceRange FullRange = SourceRange(LHS->getBeginLoc(), RHS->getEndLoc());
1340 SemaRef.Diag(LHS->getBeginLoc(), diag::note_function_suggestion)
1341 << NewFnName << FixItHint::CreateReplacement(FullRange, OS.str());
1342}
1343
1344std::pair<IdentifierInfo *, bool>
1346 llvm::hash_code Hash = llvm::hash_value(Signature);
1347 std::string IdStr = "__hlsl_rootsig_decl_" + std::to_string(Hash);
1348 IdentifierInfo *DeclIdent = &(getASTContext().Idents.get(IdStr));
1349
1350 // Check if we have already found a decl of the same name.
1351 LookupResult R(SemaRef, DeclIdent, SourceLocation(),
1353 bool Found = SemaRef.LookupQualifiedName(R, SemaRef.CurContext);
1354 return {DeclIdent, Found};
1355}
1356
1358 SourceLocation Loc, IdentifierInfo *DeclIdent,
1360
1361 if (handleRootSignatureElements(RootElements))
1362 return;
1363
1365 for (auto &RootSigElement : RootElements)
1366 Elements.push_back(RootSigElement.getElement());
1367
1368 auto *SignatureDecl = HLSLRootSignatureDecl::Create(
1369 SemaRef.getASTContext(), /*DeclContext=*/SemaRef.CurContext, Loc,
1370 DeclIdent, SemaRef.getLangOpts().HLSLRootSigVer, Elements);
1371
1372 SignatureDecl->setImplicit();
1373 SemaRef.PushOnScopeChains(SignatureDecl, SemaRef.getCurScope());
1374}
1375
1378 if (RootSigOverrideIdent) {
1379 LookupResult R(SemaRef, RootSigOverrideIdent, SourceLocation(),
1381 if (SemaRef.LookupQualifiedName(R, DC))
1382 return dyn_cast<HLSLRootSignatureDecl>(R.getFoundDecl());
1383 }
1384
1385 return nullptr;
1386}
1387
1388namespace {
1389
1390struct PerVisibilityBindingChecker {
1391 SemaHLSL *S;
1392 // We need one builder per `llvm::dxbc::ShaderVisibility` value.
1393 std::array<llvm::hlsl::BindingInfoBuilder, 8> Builders;
1394
1395 struct ElemInfo {
1396 const hlsl::RootSignatureElement *Elem;
1397 llvm::dxbc::ShaderVisibility Vis;
1398 bool Diagnosed;
1399 };
1400 llvm::SmallVector<ElemInfo> ElemInfoMap;
1401
1402 PerVisibilityBindingChecker(SemaHLSL *S) : S(S) {}
1403
1404 void trackBinding(llvm::dxbc::ShaderVisibility Visibility,
1405 llvm::dxil::ResourceClass RC, uint32_t Space,
1406 uint32_t LowerBound, uint32_t UpperBound,
1407 const hlsl::RootSignatureElement *Elem) {
1408 uint32_t BuilderIndex = llvm::to_underlying(Visibility);
1409 assert(BuilderIndex < Builders.size() &&
1410 "Not enough builders for visibility type");
1411 Builders[BuilderIndex].trackBinding(RC, Space, LowerBound, UpperBound,
1412 static_cast<const void *>(Elem));
1413
1414 static_assert(llvm::to_underlying(llvm::dxbc::ShaderVisibility::All) == 0,
1415 "'All' visibility must come first");
1416 if (Visibility == llvm::dxbc::ShaderVisibility::All)
1417 for (size_t I = 1, E = Builders.size(); I < E; ++I)
1418 Builders[I].trackBinding(RC, Space, LowerBound, UpperBound,
1419 static_cast<const void *>(Elem));
1420
1421 ElemInfoMap.push_back({Elem, Visibility, false});
1422 }
1423
1424 ElemInfo &getInfo(const hlsl::RootSignatureElement *Elem) {
1425 auto It = llvm::lower_bound(
1426 ElemInfoMap, Elem,
1427 [](const auto &LHS, const auto &RHS) { return LHS.Elem < RHS; });
1428 assert(It->Elem == Elem && "Element not in map");
1429 return *It;
1430 }
1431
1432 bool checkOverlap() {
1433 llvm::sort(ElemInfoMap, [](const auto &LHS, const auto &RHS) {
1434 return LHS.Elem < RHS.Elem;
1435 });
1436
1437 bool HadOverlap = false;
1438
1439 using llvm::hlsl::BindingInfoBuilder;
1440 auto ReportOverlap = [this,
1441 &HadOverlap](const BindingInfoBuilder &Builder,
1442 const llvm::hlsl::Binding &Reported) {
1443 HadOverlap = true;
1444
1445 const auto *Elem =
1446 static_cast<const hlsl::RootSignatureElement *>(Reported.Cookie);
1447 const llvm::hlsl::Binding &Previous = Builder.findOverlapping(Reported);
1448 const auto *PrevElem =
1449 static_cast<const hlsl::RootSignatureElement *>(Previous.Cookie);
1450
1451 ElemInfo &Info = getInfo(Elem);
1452 // We will have already diagnosed this binding if there's overlap in the
1453 // "All" visibility as well as any particular visibility.
1454 if (Info.Diagnosed)
1455 return;
1456 Info.Diagnosed = true;
1457
1458 ElemInfo &PrevInfo = getInfo(PrevElem);
1459 llvm::dxbc::ShaderVisibility CommonVis =
1460 Info.Vis == llvm::dxbc::ShaderVisibility::All ? PrevInfo.Vis
1461 : Info.Vis;
1462
1463 this->S->Diag(Elem->getLocation(), diag::err_hlsl_resource_range_overlap)
1464 << llvm::to_underlying(Reported.RC) << Reported.LowerBound
1465 << Reported.isUnbounded() << Reported.UpperBound
1466 << llvm::to_underlying(Previous.RC) << Previous.LowerBound
1467 << Previous.isUnbounded() << Previous.UpperBound << Reported.Space
1468 << CommonVis;
1469
1470 this->S->Diag(PrevElem->getLocation(),
1471 diag::note_hlsl_resource_range_here);
1472 };
1473
1474 for (BindingInfoBuilder &Builder : Builders)
1475 Builder.calculateBindingInfo(ReportOverlap);
1476
1477 return HadOverlap;
1478 }
1479};
1480
1481static CXXMethodDecl *lookupMethod(Sema &S, CXXRecordDecl *RecordDecl,
1482 StringRef Name, SourceLocation Loc) {
1483 DeclarationName DeclName(&S.getASTContext().Idents.get(Name));
1484 LookupResult Result(S, DeclName, Loc, Sema::LookupMemberName);
1485 if (!S.LookupQualifiedName(Result, static_cast<DeclContext *>(RecordDecl)))
1486 return nullptr;
1487 return cast<CXXMethodDecl>(Result.getFoundDecl());
1488}
1489
1490} // end anonymous namespace
1491
1492static bool hasCounterHandle(const CXXRecordDecl *RD) {
1493 if (RD->field_empty())
1494 return false;
1495 auto It = std::next(RD->field_begin());
1496 if (It == RD->field_end())
1497 return false;
1498 const FieldDecl *SecondField = *It;
1499 if (const auto *ResTy =
1500 SecondField->getType()->getAs<HLSLAttributedResourceType>()) {
1501 return ResTy->getAttrs().IsCounter;
1502 }
1503 return false;
1504}
1505
1508 // Define some common error handling functions
1509 bool HadError = false;
1510 auto ReportError = [this, &HadError](SourceLocation Loc, uint32_t LowerBound,
1511 uint32_t UpperBound) {
1512 HadError = true;
1513 this->Diag(Loc, diag::err_hlsl_invalid_rootsig_value)
1514 << LowerBound << UpperBound;
1515 };
1516
1517 auto ReportFloatError = [this, &HadError](SourceLocation Loc,
1518 float LowerBound,
1519 float UpperBound) {
1520 HadError = true;
1521 this->Diag(Loc, diag::err_hlsl_invalid_rootsig_value)
1522 << llvm::formatv("{0:f}", LowerBound).sstr<6>()
1523 << llvm::formatv("{0:f}", UpperBound).sstr<6>();
1524 };
1525
1526 auto VerifyRegister = [ReportError](SourceLocation Loc, uint32_t Register) {
1527 if (!llvm::hlsl::rootsig::verifyRegisterValue(Register))
1528 ReportError(Loc, 0, 0xfffffffe);
1529 };
1530
1531 auto VerifySpace = [ReportError](SourceLocation Loc, uint32_t Space) {
1532 if (!llvm::hlsl::rootsig::verifyRegisterSpace(Space))
1533 ReportError(Loc, 0, 0xffffffef);
1534 };
1535
1536 const uint32_t Version =
1537 llvm::to_underlying(SemaRef.getLangOpts().HLSLRootSigVer);
1538 const uint32_t VersionEnum = Version - 1;
1539 auto ReportFlagError = [this, &HadError, VersionEnum](SourceLocation Loc) {
1540 HadError = true;
1541 this->Diag(Loc, diag::err_hlsl_invalid_rootsig_flag)
1542 << /*version minor*/ VersionEnum;
1543 };
1544
1545 // Iterate through the elements and do basic validations
1546 for (const hlsl::RootSignatureElement &RootSigElem : Elements) {
1547 SourceLocation Loc = RootSigElem.getLocation();
1548 const llvm::hlsl::rootsig::RootElement &Elem = RootSigElem.getElement();
1549 if (const auto *Descriptor =
1550 std::get_if<llvm::hlsl::rootsig::RootDescriptor>(&Elem)) {
1551 VerifyRegister(Loc, Descriptor->Reg.Number);
1552 VerifySpace(Loc, Descriptor->Space);
1553
1554 if (!llvm::hlsl::rootsig::verifyRootDescriptorFlag(Version,
1555 Descriptor->Flags))
1556 ReportFlagError(Loc);
1557 } else if (const auto *Constants =
1558 std::get_if<llvm::hlsl::rootsig::RootConstants>(&Elem)) {
1559 VerifyRegister(Loc, Constants->Reg.Number);
1560 VerifySpace(Loc, Constants->Space);
1561 } else if (const auto *Sampler =
1562 std::get_if<llvm::hlsl::rootsig::StaticSampler>(&Elem)) {
1563 VerifyRegister(Loc, Sampler->Reg.Number);
1564 VerifySpace(Loc, Sampler->Space);
1565
1566 assert(!std::isnan(Sampler->MaxLOD) && !std::isnan(Sampler->MinLOD) &&
1567 "By construction, parseFloatParam can't produce a NaN from a "
1568 "float_literal token");
1569
1570 if (!llvm::hlsl::rootsig::verifyMaxAnisotropy(Sampler->MaxAnisotropy))
1571 ReportError(Loc, 0, 16);
1572 if (!llvm::hlsl::rootsig::verifyMipLODBias(Sampler->MipLODBias))
1573 ReportFloatError(Loc, -16.f, 15.99f);
1574 } else if (const auto *Clause =
1575 std::get_if<llvm::hlsl::rootsig::DescriptorTableClause>(
1576 &Elem)) {
1577 VerifyRegister(Loc, Clause->Reg.Number);
1578 VerifySpace(Loc, Clause->Space);
1579
1580 if (!llvm::hlsl::rootsig::verifyNumDescriptors(Clause->NumDescriptors)) {
1581 // NumDescriptor could techincally be ~0u but that is reserved for
1582 // unbounded, so the diagnostic will not report that as a valid int
1583 // value
1584 ReportError(Loc, 1, 0xfffffffe);
1585 }
1586
1587 if (!llvm::hlsl::rootsig::verifyDescriptorRangeFlag(Version, Clause->Type,
1588 Clause->Flags))
1589 ReportFlagError(Loc);
1590 }
1591 }
1592
1593 PerVisibilityBindingChecker BindingChecker(this);
1594 SmallVector<std::pair<const llvm::hlsl::rootsig::DescriptorTableClause *,
1596 UnboundClauses;
1597
1598 for (const hlsl::RootSignatureElement &RootSigElem : Elements) {
1599 const llvm::hlsl::rootsig::RootElement &Elem = RootSigElem.getElement();
1600 if (const auto *Descriptor =
1601 std::get_if<llvm::hlsl::rootsig::RootDescriptor>(&Elem)) {
1602 uint32_t LowerBound(Descriptor->Reg.Number);
1603 uint32_t UpperBound(LowerBound); // inclusive range
1604
1605 BindingChecker.trackBinding(
1606 Descriptor->Visibility,
1607 static_cast<llvm::dxil::ResourceClass>(Descriptor->Type),
1608 Descriptor->Space, LowerBound, UpperBound, &RootSigElem);
1609 } else if (const auto *Constants =
1610 std::get_if<llvm::hlsl::rootsig::RootConstants>(&Elem)) {
1611 uint32_t LowerBound(Constants->Reg.Number);
1612 uint32_t UpperBound(LowerBound); // inclusive range
1613
1614 BindingChecker.trackBinding(
1615 Constants->Visibility, llvm::dxil::ResourceClass::CBuffer,
1616 Constants->Space, LowerBound, UpperBound, &RootSigElem);
1617 } else if (const auto *Sampler =
1618 std::get_if<llvm::hlsl::rootsig::StaticSampler>(&Elem)) {
1619 uint32_t LowerBound(Sampler->Reg.Number);
1620 uint32_t UpperBound(LowerBound); // inclusive range
1621
1622 BindingChecker.trackBinding(
1623 Sampler->Visibility, llvm::dxil::ResourceClass::Sampler,
1624 Sampler->Space, LowerBound, UpperBound, &RootSigElem);
1625 } else if (const auto *Clause =
1626 std::get_if<llvm::hlsl::rootsig::DescriptorTableClause>(
1627 &Elem)) {
1628 // We'll process these once we see the table element.
1629 UnboundClauses.emplace_back(Clause, &RootSigElem);
1630 } else if (const auto *Table =
1631 std::get_if<llvm::hlsl::rootsig::DescriptorTable>(&Elem)) {
1632 assert(UnboundClauses.size() == Table->NumClauses &&
1633 "Number of unbound elements must match the number of clauses");
1634 bool HasAnySampler = false;
1635 bool HasAnyNonSampler = false;
1636 uint64_t Offset = 0;
1637 bool IsPrevUnbound = false;
1638 for (const auto &[Clause, ClauseElem] : UnboundClauses) {
1639 SourceLocation Loc = ClauseElem->getLocation();
1640 if (Clause->Type == llvm::dxil::ResourceClass::Sampler)
1641 HasAnySampler = true;
1642 else
1643 HasAnyNonSampler = true;
1644
1645 if (HasAnySampler && HasAnyNonSampler)
1646 Diag(Loc, diag::err_hlsl_invalid_mixed_resources);
1647
1648 // Relevant error will have already been reported above and needs to be
1649 // fixed before we can conduct further analysis, so shortcut error
1650 // return
1651 if (Clause->NumDescriptors == 0)
1652 return true;
1653
1654 bool IsAppending =
1655 Clause->Offset == llvm::hlsl::rootsig::DescriptorTableOffsetAppend;
1656 if (!IsAppending)
1657 Offset = Clause->Offset;
1658
1659 uint64_t RangeBound = llvm::hlsl::rootsig::computeRangeBound(
1660 Offset, Clause->NumDescriptors);
1661
1662 if (IsPrevUnbound && IsAppending)
1663 Diag(Loc, diag::err_hlsl_appending_onto_unbound);
1664 else if (!llvm::hlsl::rootsig::verifyNoOverflowedOffset(RangeBound))
1665 Diag(Loc, diag::err_hlsl_offset_overflow) << Offset << RangeBound;
1666
1667 // Update offset to be 1 past this range's bound
1668 Offset = RangeBound + 1;
1669 IsPrevUnbound = Clause->NumDescriptors ==
1670 llvm::hlsl::rootsig::NumDescriptorsUnbounded;
1671
1672 // Compute the register bounds and track resource binding
1673 uint32_t LowerBound(Clause->Reg.Number);
1674 uint32_t UpperBound = llvm::hlsl::rootsig::computeRangeBound(
1675 LowerBound, Clause->NumDescriptors);
1676
1677 BindingChecker.trackBinding(
1678 Table->Visibility,
1679 static_cast<llvm::dxil::ResourceClass>(Clause->Type), Clause->Space,
1680 LowerBound, UpperBound, ClauseElem);
1681 }
1682 UnboundClauses.clear();
1683 }
1684 }
1685
1686 return BindingChecker.checkOverlap();
1687}
1688
1690 if (AL.getNumArgs() != 1) {
1691 Diag(AL.getLoc(), diag::err_attribute_wrong_number_arguments) << AL << 1;
1692 return;
1693 }
1694
1696 if (auto *RS = D->getAttr<RootSignatureAttr>()) {
1697 if (RS->getSignatureIdent() != Ident) {
1698 Diag(AL.getLoc(), diag::err_disallowed_duplicate_attribute) << RS;
1699 return;
1700 }
1701
1702 Diag(AL.getLoc(), diag::warn_duplicate_attribute_exact) << RS;
1703 return;
1704 }
1705
1707 if (SemaRef.LookupQualifiedName(R, D->getDeclContext()))
1708 if (auto *SignatureDecl =
1709 dyn_cast<HLSLRootSignatureDecl>(R.getFoundDecl())) {
1710 D->addAttr(::new (getASTContext()) RootSignatureAttr(
1711 getASTContext(), AL, Ident, SignatureDecl));
1712 }
1713}
1714
1716 llvm::VersionTuple SMVersion =
1717 getASTContext().getTargetInfo().getTriple().getOSVersion();
1718 bool IsDXIL = getASTContext().getTargetInfo().getTriple().getArch() ==
1719 llvm::Triple::dxil;
1720
1721 uint32_t ZMax = 1024;
1722 uint32_t ThreadMax = 1024;
1723 if (IsDXIL && SMVersion.getMajor() <= 4) {
1724 ZMax = 1;
1725 ThreadMax = 768;
1726 } else if (IsDXIL && SMVersion.getMajor() == 5) {
1727 ZMax = 64;
1728 ThreadMax = 1024;
1729 }
1730
1731 uint32_t X;
1732 if (!SemaRef.checkUInt32Argument(AL, AL.getArgAsExpr(0), X))
1733 return;
1734 if (X > 1024) {
1735 Diag(AL.getArgAsExpr(0)->getExprLoc(),
1736 diag::err_hlsl_numthreads_argument_oor)
1737 << 0 << 1024;
1738 return;
1739 }
1740 uint32_t Y;
1741 if (!SemaRef.checkUInt32Argument(AL, AL.getArgAsExpr(1), Y))
1742 return;
1743 if (Y > 1024) {
1744 Diag(AL.getArgAsExpr(1)->getExprLoc(),
1745 diag::err_hlsl_numthreads_argument_oor)
1746 << 1 << 1024;
1747 return;
1748 }
1749 uint32_t Z;
1750 if (!SemaRef.checkUInt32Argument(AL, AL.getArgAsExpr(2), Z))
1751 return;
1752 if (Z > ZMax) {
1753 SemaRef.Diag(AL.getArgAsExpr(2)->getExprLoc(),
1754 diag::err_hlsl_numthreads_argument_oor)
1755 << 2 << ZMax;
1756 return;
1757 }
1758
1759 if (X * Y * Z > ThreadMax) {
1760 Diag(AL.getLoc(), diag::err_hlsl_numthreads_invalid) << ThreadMax;
1761 return;
1762 }
1763
1764 HLSLNumThreadsAttr *NewAttr = mergeNumThreadsAttr(D, AL, X, Y, Z);
1765 if (NewAttr)
1766 D->addAttr(NewAttr);
1767}
1768
1769static bool isValidWaveSizeValue(unsigned Value) {
1770 return llvm::isPowerOf2_32(Value) && Value >= 4 && Value <= 128;
1771}
1772
1774 // validate that the wavesize argument is a power of 2 between 4 and 128
1775 // inclusive
1776 unsigned SpelledArgsCount = AL.getNumArgs();
1777 if (SpelledArgsCount == 0 || SpelledArgsCount > 3)
1778 return;
1779
1780 uint32_t Min;
1781 if (!SemaRef.checkUInt32Argument(AL, AL.getArgAsExpr(0), Min))
1782 return;
1783
1784 uint32_t Max = 0;
1785 if (SpelledArgsCount > 1 &&
1786 !SemaRef.checkUInt32Argument(AL, AL.getArgAsExpr(1), Max))
1787 return;
1788
1789 uint32_t Preferred = 0;
1790 if (SpelledArgsCount > 2 &&
1791 !SemaRef.checkUInt32Argument(AL, AL.getArgAsExpr(2), Preferred))
1792 return;
1793
1794 if (SpelledArgsCount > 2) {
1795 if (!isValidWaveSizeValue(Preferred)) {
1796 Diag(AL.getArgAsExpr(2)->getExprLoc(),
1797 diag::err_attribute_power_of_two_in_range)
1798 << AL << llvm::dxil::MinWaveSize << llvm::dxil::MaxWaveSize
1799 << Preferred;
1800 return;
1801 }
1802 // Preferred not in range.
1803 if (Preferred < Min || Preferred > Max) {
1804 Diag(AL.getArgAsExpr(2)->getExprLoc(),
1805 diag::err_attribute_power_of_two_in_range)
1806 << AL << Min << Max << Preferred;
1807 return;
1808 }
1809 } else if (SpelledArgsCount > 1) {
1810 if (!isValidWaveSizeValue(Max)) {
1811 Diag(AL.getArgAsExpr(1)->getExprLoc(),
1812 diag::err_attribute_power_of_two_in_range)
1813 << AL << llvm::dxil::MinWaveSize << llvm::dxil::MaxWaveSize << Max;
1814 return;
1815 }
1816 if (Max < Min) {
1817 Diag(AL.getLoc(), diag::err_attribute_argument_invalid) << AL << 1;
1818 return;
1819 } else if (Max == Min) {
1820 Diag(AL.getLoc(), diag::warn_attr_min_eq_max) << AL;
1821 }
1822 } else {
1823 if (!isValidWaveSizeValue(Min)) {
1824 Diag(AL.getArgAsExpr(0)->getExprLoc(),
1825 diag::err_attribute_power_of_two_in_range)
1826 << AL << llvm::dxil::MinWaveSize << llvm::dxil::MaxWaveSize << Min;
1827 return;
1828 }
1829 }
1830
1831 HLSLWaveSizeAttr *NewAttr =
1832 mergeWaveSizeAttr(D, AL, Min, Max, Preferred, SpelledArgsCount);
1833 if (NewAttr)
1834 D->addAttr(NewAttr);
1835}
1836
1838 uint32_t ID;
1839 if (!SemaRef.checkUInt32Argument(AL, AL.getArgAsExpr(0), ID))
1840 return;
1841 D->addAttr(::new (getASTContext())
1842 HLSLVkExtBuiltinInputAttr(getASTContext(), AL, ID));
1843}
1844
1846 D->addAttr(::new (getASTContext())
1847 HLSLVkPushConstantAttr(getASTContext(), AL));
1848}
1849
1851 uint32_t Id;
1852 if (!SemaRef.checkUInt32Argument(AL, AL.getArgAsExpr(0), Id))
1853 return;
1854 HLSLVkConstantIdAttr *NewAttr = mergeVkConstantIdAttr(D, AL, Id);
1855 if (NewAttr)
1856 D->addAttr(NewAttr);
1857}
1858
1860 uint32_t Binding = 0;
1861 if (!SemaRef.checkUInt32Argument(AL, AL.getArgAsExpr(0), Binding))
1862 return;
1863 uint32_t Set = 0;
1864 if (AL.getNumArgs() > 1 &&
1865 !SemaRef.checkUInt32Argument(AL, AL.getArgAsExpr(1), Set))
1866 return;
1867
1868 D->addAttr(::new (getASTContext())
1869 HLSLVkBindingAttr(getASTContext(), AL, Binding, Set));
1870}
1871
1873 uint32_t Location;
1874 if (!SemaRef.checkUInt32Argument(AL, AL.getArgAsExpr(0), Location))
1875 return;
1876
1877 D->addAttr(::new (getASTContext())
1878 HLSLVkLocationAttr(getASTContext(), AL, Location));
1879}
1880
1882 const auto *VT = T->getAs<VectorType>();
1883
1884 if (!T->hasUnsignedIntegerRepresentation() ||
1885 (VT && VT->getNumElements() > 3)) {
1886 Diag(AL.getLoc(), diag::err_hlsl_attr_invalid_type)
1887 << AL << "uint/uint2/uint3";
1888 return false;
1889 }
1890
1891 return true;
1892}
1893
1895 const auto *VT = T->getAs<VectorType>();
1896 if (!T->hasFloatingRepresentation() || (VT && VT->getNumElements() > 4)) {
1897 Diag(AL.getLoc(), diag::err_hlsl_attr_invalid_type)
1898 << AL << "float/float1/float2/float3/float4";
1899 return false;
1900 }
1901
1902 return true;
1903}
1904
1906 std::optional<unsigned> Index) {
1907 std::string SemanticName = AL.getAttrName()->getName().upper();
1908
1909 auto *VD = cast<ValueDecl>(D);
1910 QualType ValueType = VD->getType();
1911 if (auto *FD = dyn_cast<FunctionDecl>(D))
1912 ValueType = FD->getReturnType();
1913
1914 bool IsOutput = false;
1915 if (HLSLParamModifierAttr *MA = D->getAttr<HLSLParamModifierAttr>()) {
1916 if (MA->isOut()) {
1917 IsOutput = true;
1918 ValueType = cast<ReferenceType>(ValueType)->getPointeeType();
1919 }
1920 }
1921
1922 if (SemanticName == "SV_DISPATCHTHREADID") {
1923 diagnoseInputIDType(ValueType, AL);
1924 if (IsOutput)
1925 Diag(AL.getLoc(), diag::err_hlsl_semantic_output_not_supported) << AL;
1926 if (Index.has_value())
1927 Diag(AL.getLoc(), diag::err_hlsl_semantic_indexing_not_supported) << AL;
1929 return;
1930 }
1931
1932 if (SemanticName == "SV_GROUPINDEX") {
1933 if (IsOutput)
1934 Diag(AL.getLoc(), diag::err_hlsl_semantic_output_not_supported) << AL;
1935 if (Index.has_value())
1936 Diag(AL.getLoc(), diag::err_hlsl_semantic_indexing_not_supported) << AL;
1938 return;
1939 }
1940
1941 if (SemanticName == "SV_GROUPTHREADID") {
1942 diagnoseInputIDType(ValueType, AL);
1943 if (IsOutput)
1944 Diag(AL.getLoc(), diag::err_hlsl_semantic_output_not_supported) << AL;
1945 if (Index.has_value())
1946 Diag(AL.getLoc(), diag::err_hlsl_semantic_indexing_not_supported) << AL;
1948 return;
1949 }
1950
1951 if (SemanticName == "SV_GROUPID") {
1952 diagnoseInputIDType(ValueType, AL);
1953 if (IsOutput)
1954 Diag(AL.getLoc(), diag::err_hlsl_semantic_output_not_supported) << AL;
1955 if (Index.has_value())
1956 Diag(AL.getLoc(), diag::err_hlsl_semantic_indexing_not_supported) << AL;
1958 return;
1959 }
1960
1961 if (SemanticName == "SV_POSITION") {
1962 const auto *VT = ValueType->getAs<VectorType>();
1963 if (!ValueType->hasFloatingRepresentation() ||
1964 (VT && VT->getNumElements() > 4))
1965 Diag(AL.getLoc(), diag::err_hlsl_attr_invalid_type)
1966 << AL << "float/float1/float2/float3/float4";
1968 return;
1969 }
1970
1971 if (SemanticName == "SV_VERTEXID") {
1972 uint64_t SizeInBits = SemaRef.Context.getTypeSize(ValueType);
1973 if (!ValueType->isUnsignedIntegerType() || SizeInBits != 32)
1974 Diag(AL.getLoc(), diag::err_hlsl_attr_invalid_type) << AL << "uint";
1976 return;
1977 }
1978
1979 if (SemanticName == "SV_TARGET") {
1980 const auto *VT = ValueType->getAs<VectorType>();
1981 if (!ValueType->hasFloatingRepresentation() ||
1982 (VT && VT->getNumElements() > 4))
1983 Diag(AL.getLoc(), diag::err_hlsl_attr_invalid_type)
1984 << AL << "float/float1/float2/float3/float4";
1986 return;
1987 }
1988
1989 Diag(AL.getLoc(), diag::err_hlsl_unknown_semantic) << AL;
1990}
1991
1993 uint32_t IndexValue(0), ExplicitIndex(0);
1994 if (!SemaRef.checkUInt32Argument(AL, AL.getArgAsExpr(0), IndexValue) ||
1995 !SemaRef.checkUInt32Argument(AL, AL.getArgAsExpr(1), ExplicitIndex)) {
1996 assert(0 && "HLSLUnparsedSemantic is expected to have 2 int arguments.");
1997 }
1998 assert(IndexValue > 0 ? ExplicitIndex : true);
1999 std::optional<unsigned> Index =
2000 ExplicitIndex ? std::optional<unsigned>(IndexValue) : std::nullopt;
2001
2002 if (AL.getAttrName()->getName().starts_with_insensitive("SV_"))
2003 diagnoseSystemSemanticAttr(D, AL, Index);
2004 else
2006}
2007
2010 Diag(AL.getLoc(), diag::err_hlsl_attr_invalid_ast_node)
2011 << AL << "shader constant in a constant buffer";
2012 return;
2013 }
2014
2015 uint32_t SubComponent;
2016 if (!SemaRef.checkUInt32Argument(AL, AL.getArgAsExpr(0), SubComponent))
2017 return;
2018 uint32_t Component;
2019 if (!SemaRef.checkUInt32Argument(AL, AL.getArgAsExpr(1), Component))
2020 return;
2021
2022 QualType T = cast<VarDecl>(D)->getType().getCanonicalType();
2023 // Check if T is an array or struct type.
2024 // TODO: mark matrix type as aggregate type.
2025 bool IsAggregateTy = (T->isArrayType() || T->isStructureType());
2026
2027 // Check Component is valid for T.
2028 if (Component) {
2029 unsigned Size = getASTContext().getTypeSize(T);
2030 if (IsAggregateTy) {
2031 Diag(AL.getLoc(), diag::err_hlsl_invalid_register_or_packoffset);
2032 return;
2033 } else {
2034 // Make sure Component + sizeof(T) <= 4.
2035 if ((Component * 32 + Size) > 128) {
2036 Diag(AL.getLoc(), diag::err_hlsl_packoffset_cross_reg_boundary);
2037 return;
2038 }
2039 QualType EltTy = T;
2040 if (const auto *VT = T->getAs<VectorType>())
2041 EltTy = VT->getElementType();
2042 unsigned Align = getASTContext().getTypeAlign(EltTy);
2043 if (Align > 32 && Component == 1) {
2044 // NOTE: Component 3 will hit err_hlsl_packoffset_cross_reg_boundary.
2045 // So we only need to check Component 1 here.
2046 Diag(AL.getLoc(), diag::err_hlsl_packoffset_alignment_mismatch)
2047 << Align << EltTy;
2048 return;
2049 }
2050 }
2051 }
2052
2053 D->addAttr(::new (getASTContext()) HLSLPackOffsetAttr(
2054 getASTContext(), AL, SubComponent, Component));
2055}
2056
2058 StringRef Str;
2059 SourceLocation ArgLoc;
2060 if (!SemaRef.checkStringLiteralArgumentAttr(AL, 0, Str, &ArgLoc))
2061 return;
2062
2063 llvm::Triple::EnvironmentType ShaderType;
2064 if (!HLSLShaderAttr::ConvertStrToEnvironmentType(Str, ShaderType)) {
2065 Diag(AL.getLoc(), diag::warn_attribute_type_not_supported)
2066 << AL << Str << ArgLoc;
2067 return;
2068 }
2069
2070 // FIXME: check function match the shader stage.
2071
2072 HLSLShaderAttr *NewAttr = mergeShaderAttr(D, AL, ShaderType);
2073 if (NewAttr)
2074 D->addAttr(NewAttr);
2075}
2076
2078 Sema &S, QualType Wrapped, ArrayRef<const Attr *> AttrList,
2079 QualType &ResType, HLSLAttributedResourceLocInfo *LocInfo) {
2080 assert(AttrList.size() && "expected list of resource attributes");
2081
2082 QualType ContainedTy = QualType();
2083 TypeSourceInfo *ContainedTyInfo = nullptr;
2084 SourceLocation LocBegin = AttrList[0]->getRange().getBegin();
2085 SourceLocation LocEnd = AttrList[0]->getRange().getEnd();
2086
2087 HLSLAttributedResourceType::Attributes ResAttrs;
2088
2089 bool HasResourceClass = false;
2090 bool HasResourceDimension = false;
2091 for (const Attr *A : AttrList) {
2092 if (!A)
2093 continue;
2094 LocEnd = A->getRange().getEnd();
2095 switch (A->getKind()) {
2096 case attr::HLSLResourceClass: {
2097 ResourceClass RC = cast<HLSLResourceClassAttr>(A)->getResourceClass();
2098 if (HasResourceClass) {
2099 S.Diag(A->getLocation(), ResAttrs.ResourceClass == RC
2100 ? diag::warn_duplicate_attribute_exact
2101 : diag::warn_duplicate_attribute)
2102 << A;
2103 return false;
2104 }
2105 ResAttrs.ResourceClass = RC;
2106 HasResourceClass = true;
2107 break;
2108 }
2109 case attr::HLSLResourceDimension: {
2110 llvm::dxil::ResourceDimension RD =
2111 cast<HLSLResourceDimensionAttr>(A)->getDimension();
2112 if (HasResourceDimension) {
2113 S.Diag(A->getLocation(), ResAttrs.ResourceDimension == RD
2114 ? diag::warn_duplicate_attribute_exact
2115 : diag::warn_duplicate_attribute)
2116 << A;
2117 return false;
2118 }
2119 ResAttrs.ResourceDimension = RD;
2120 HasResourceDimension = true;
2121 break;
2122 }
2123 case attr::HLSLROV:
2124 if (ResAttrs.IsROV) {
2125 S.Diag(A->getLocation(), diag::warn_duplicate_attribute_exact) << A;
2126 return false;
2127 }
2128 ResAttrs.IsROV = true;
2129 break;
2130 case attr::HLSLRawBuffer:
2131 if (ResAttrs.RawBuffer) {
2132 S.Diag(A->getLocation(), diag::warn_duplicate_attribute_exact) << A;
2133 return false;
2134 }
2135 ResAttrs.RawBuffer = true;
2136 break;
2137 case attr::HLSLIsCounter:
2138 if (ResAttrs.IsCounter) {
2139 S.Diag(A->getLocation(), diag::warn_duplicate_attribute_exact) << A;
2140 return false;
2141 }
2142 ResAttrs.IsCounter = true;
2143 break;
2144 case attr::HLSLContainedType: {
2145 const HLSLContainedTypeAttr *CTAttr = cast<HLSLContainedTypeAttr>(A);
2146 QualType Ty = CTAttr->getType();
2147 if (!ContainedTy.isNull()) {
2148 S.Diag(A->getLocation(), ContainedTy == Ty
2149 ? diag::warn_duplicate_attribute_exact
2150 : diag::warn_duplicate_attribute)
2151 << A;
2152 return false;
2153 }
2154 ContainedTy = Ty;
2155 ContainedTyInfo = CTAttr->getTypeLoc();
2156 break;
2157 }
2158 default:
2159 llvm_unreachable("unhandled resource attribute type");
2160 }
2161 }
2162
2163 if (!HasResourceClass) {
2164 S.Diag(AttrList.back()->getRange().getEnd(),
2165 diag::err_hlsl_missing_resource_class);
2166 return false;
2167 }
2168
2170 Wrapped, ContainedTy, ResAttrs);
2171
2172 if (LocInfo && ContainedTyInfo) {
2173 LocInfo->Range = SourceRange(LocBegin, LocEnd);
2174 LocInfo->ContainedTyInfo = ContainedTyInfo;
2175 }
2176 return true;
2177}
2178
2179// Validates and creates an HLSL attribute that is applied as type attribute on
2180// HLSL resource. The attributes are collected in HLSLResourcesTypeAttrs and at
2181// the end of the declaration they are applied to the declaration type by
2182// wrapping it in HLSLAttributedResourceType.
2184 // only allow resource type attributes on intangible types
2185 if (!T->isHLSLResourceType()) {
2186 Diag(AL.getLoc(), diag::err_hlsl_attribute_needs_intangible_type)
2187 << AL << getASTContext().HLSLResourceTy;
2188 return false;
2189 }
2190
2191 // validate number of arguments
2192 if (!AL.checkExactlyNumArgs(SemaRef, AL.getMinArgs()))
2193 return false;
2194
2195 Attr *A = nullptr;
2196
2200 {
2201 AttributeCommonInfo::AS_CXX11, 0, false /*IsAlignas*/,
2202 false /*IsRegularKeywordAttribute*/
2203 });
2204
2205 switch (AL.getKind()) {
2206 case ParsedAttr::AT_HLSLResourceClass: {
2207 if (!AL.isArgIdent(0)) {
2208 Diag(AL.getLoc(), diag::err_attribute_argument_type)
2209 << AL << AANT_ArgumentIdentifier;
2210 return false;
2211 }
2212
2213 IdentifierLoc *Loc = AL.getArgAsIdent(0);
2214 StringRef Identifier = Loc->getIdentifierInfo()->getName();
2215 SourceLocation ArgLoc = Loc->getLoc();
2216
2217 // Validate resource class value
2218 ResourceClass RC;
2219 if (!HLSLResourceClassAttr::ConvertStrToResourceClass(Identifier, RC)) {
2220 Diag(ArgLoc, diag::warn_attribute_type_not_supported)
2221 << "ResourceClass" << Identifier;
2222 return false;
2223 }
2224 A = HLSLResourceClassAttr::Create(getASTContext(), RC, ACI);
2225 break;
2226 }
2227
2228 case ParsedAttr::AT_HLSLResourceDimension: {
2229 StringRef Identifier;
2230 SourceLocation ArgLoc;
2231 if (!SemaRef.checkStringLiteralArgumentAttr(AL, 0, Identifier, &ArgLoc))
2232 return false;
2233
2234 // Validate resource dimension value
2235 llvm::dxil::ResourceDimension RD;
2236 if (!HLSLResourceDimensionAttr::ConvertStrToResourceDimension(Identifier,
2237 RD)) {
2238 Diag(ArgLoc, diag::warn_attribute_type_not_supported)
2239 << "ResourceDimension" << Identifier;
2240 return false;
2241 }
2242 A = HLSLResourceDimensionAttr::Create(getASTContext(), RD, ACI);
2243 break;
2244 }
2245
2246 case ParsedAttr::AT_HLSLROV:
2247 A = HLSLROVAttr::Create(getASTContext(), ACI);
2248 break;
2249
2250 case ParsedAttr::AT_HLSLRawBuffer:
2251 A = HLSLRawBufferAttr::Create(getASTContext(), ACI);
2252 break;
2253
2254 case ParsedAttr::AT_HLSLIsCounter:
2255 A = HLSLIsCounterAttr::Create(getASTContext(), ACI);
2256 break;
2257
2258 case ParsedAttr::AT_HLSLContainedType: {
2259 if (AL.getNumArgs() != 1 && !AL.hasParsedType()) {
2260 Diag(AL.getLoc(), diag::err_attribute_wrong_number_arguments) << AL << 1;
2261 return false;
2262 }
2263
2264 TypeSourceInfo *TSI = nullptr;
2265 QualType QT = SemaRef.GetTypeFromParser(AL.getTypeArg(), &TSI);
2266 assert(TSI && "no type source info for attribute argument");
2267 if (SemaRef.RequireCompleteType(TSI->getTypeLoc().getBeginLoc(), QT,
2268 diag::err_incomplete_type))
2269 return false;
2270 A = HLSLContainedTypeAttr::Create(getASTContext(), TSI, ACI);
2271 break;
2272 }
2273
2274 default:
2275 llvm_unreachable("unhandled HLSL attribute");
2276 }
2277
2278 HLSLResourcesTypeAttrs.emplace_back(A);
2279 return true;
2280}
2281
2282// Combines all resource type attributes and creates HLSLAttributedResourceType.
2284 if (!HLSLResourcesTypeAttrs.size())
2285 return CurrentType;
2286
2287 QualType QT = CurrentType;
2290 HLSLResourcesTypeAttrs, QT, &LocInfo)) {
2291 const HLSLAttributedResourceType *RT =
2293
2294 // Temporarily store TypeLoc information for the new type.
2295 // It will be transferred to HLSLAttributesResourceTypeLoc
2296 // shortly after the type is created by TypeSpecLocFiller which
2297 // will call the TakeLocForHLSLAttribute method below.
2298 LocsForHLSLAttributedResources.insert(std::pair(RT, LocInfo));
2299 }
2300 HLSLResourcesTypeAttrs.clear();
2301 return QT;
2302}
2303
2304// Returns source location for the HLSLAttributedResourceType
2306SemaHLSL::TakeLocForHLSLAttribute(const HLSLAttributedResourceType *RT) {
2307 HLSLAttributedResourceLocInfo LocInfo = {};
2308 auto I = LocsForHLSLAttributedResources.find(RT);
2309 if (I != LocsForHLSLAttributedResources.end()) {
2310 LocInfo = I->second;
2311 LocsForHLSLAttributedResources.erase(I);
2312 return LocInfo;
2313 }
2314 LocInfo.Range = SourceRange();
2315 return LocInfo;
2316}
2317
2318// Walks though the global variable declaration, collects all resource binding
2319// requirements and adds them to Bindings
2320void SemaHLSL::collectResourceBindingsOnUserRecordDecl(const VarDecl *VD,
2321 const RecordType *RT) {
2322 const RecordDecl *RD = RT->getDecl()->getDefinitionOrSelf();
2323 for (FieldDecl *FD : RD->fields()) {
2324 const Type *Ty = FD->getType()->getUnqualifiedDesugaredType();
2325
2326 // Unwrap arrays
2327 // FIXME: Calculate array size while unwrapping
2328 assert(!Ty->isIncompleteArrayType() &&
2329 "incomplete arrays inside user defined types are not supported");
2330 while (Ty->isConstantArrayType()) {
2333 }
2334
2335 if (!Ty->isRecordType())
2336 continue;
2337
2338 if (const HLSLAttributedResourceType *AttrResType =
2339 HLSLAttributedResourceType::findHandleTypeOnResource(Ty)) {
2340 // Add a new DeclBindingInfo to Bindings if it does not already exist
2341 ResourceClass RC = AttrResType->getAttrs().ResourceClass;
2342 DeclBindingInfo *DBI = Bindings.getDeclBindingInfo(VD, RC);
2343 if (!DBI)
2344 Bindings.addDeclBindingInfo(VD, RC);
2345 } else if (const RecordType *RT = dyn_cast<RecordType>(Ty)) {
2346 // Recursively scan embedded struct or class; it would be nice to do this
2347 // without recursion, but tricky to correctly calculate the size of the
2348 // binding, which is something we are probably going to need to do later
2349 // on. Hopefully nesting of structs in structs too many levels is
2350 // unlikely.
2351 collectResourceBindingsOnUserRecordDecl(VD, RT);
2352 }
2353 }
2354}
2355
2356// Diagnose localized register binding errors for a single binding; does not
2357// diagnose resource binding on user record types, that will be done later
2358// in processResourceBindingOnDecl based on the information collected in
2359// collectResourceBindingsOnVarDecl.
2360// Returns false if the register binding is not valid.
2362 Decl *D, RegisterType RegType,
2363 bool SpecifiedSpace) {
2364 int RegTypeNum = static_cast<int>(RegType);
2365
2366 // check if the decl type is groupshared
2367 if (D->hasAttr<HLSLGroupSharedAddressSpaceAttr>()) {
2368 S.Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
2369 return false;
2370 }
2371
2372 // Cbuffers and Tbuffers are HLSLBufferDecl types
2373 if (HLSLBufferDecl *CBufferOrTBuffer = dyn_cast<HLSLBufferDecl>(D)) {
2374 ResourceClass RC = CBufferOrTBuffer->isCBuffer() ? ResourceClass::CBuffer
2375 : ResourceClass::SRV;
2376 if (RegType == getRegisterType(RC))
2377 return true;
2378
2379 S.Diag(D->getLocation(), diag::err_hlsl_binding_type_mismatch)
2380 << RegTypeNum;
2381 return false;
2382 }
2383
2384 // Samplers, UAVs, and SRVs are VarDecl types
2385 assert(isa<VarDecl>(D) && "D is expected to be VarDecl or HLSLBufferDecl");
2386 VarDecl *VD = cast<VarDecl>(D);
2387
2388 // Resource
2389 if (const HLSLAttributedResourceType *AttrResType =
2390 HLSLAttributedResourceType::findHandleTypeOnResource(
2391 VD->getType().getTypePtr())) {
2392 if (RegType == getRegisterType(AttrResType))
2393 return true;
2394
2395 S.Diag(D->getLocation(), diag::err_hlsl_binding_type_mismatch)
2396 << RegTypeNum;
2397 return false;
2398 }
2399
2400 const clang::Type *Ty = VD->getType().getTypePtr();
2401 while (Ty->isArrayType())
2403
2404 // Basic types
2405 if (Ty->isArithmeticType() || Ty->isVectorType()) {
2406 bool DeclaredInCOrTBuffer = isa<HLSLBufferDecl>(D->getDeclContext());
2407 if (SpecifiedSpace && !DeclaredInCOrTBuffer)
2408 S.Diag(ArgLoc, diag::err_hlsl_space_on_global_constant);
2409
2410 if (!DeclaredInCOrTBuffer && (Ty->isIntegralType(S.getASTContext()) ||
2411 Ty->isFloatingType() || Ty->isVectorType())) {
2412 // Register annotation on default constant buffer declaration ($Globals)
2413 if (RegType == RegisterType::CBuffer)
2414 S.Diag(ArgLoc, diag::warn_hlsl_deprecated_register_type_b);
2415 else if (RegType != RegisterType::C)
2416 S.Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
2417 else
2418 return true;
2419 } else {
2420 if (RegType == RegisterType::C)
2421 S.Diag(ArgLoc, diag::warn_hlsl_register_type_c_packoffset);
2422 else
2423 S.Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
2424 }
2425 return false;
2426 }
2427 if (Ty->isRecordType())
2428 // RecordTypes will be diagnosed in processResourceBindingOnDecl
2429 // that is called from ActOnVariableDeclarator
2430 return true;
2431
2432 // Anything else is an error
2433 S.Diag(ArgLoc, diag::err_hlsl_binding_type_mismatch) << RegTypeNum;
2434 return false;
2435}
2436
2438 RegisterType regType) {
2439 // make sure that there are no two register annotations
2440 // applied to the decl with the same register type
2441 bool RegisterTypesDetected[5] = {false};
2442 RegisterTypesDetected[static_cast<int>(regType)] = true;
2443
2444 for (auto it = TheDecl->attr_begin(); it != TheDecl->attr_end(); ++it) {
2445 if (HLSLResourceBindingAttr *attr =
2446 dyn_cast<HLSLResourceBindingAttr>(*it)) {
2447
2448 RegisterType otherRegType = attr->getRegisterType();
2449 if (RegisterTypesDetected[static_cast<int>(otherRegType)]) {
2450 int otherRegTypeNum = static_cast<int>(otherRegType);
2451 S.Diag(TheDecl->getLocation(),
2452 diag::err_hlsl_duplicate_register_annotation)
2453 << otherRegTypeNum;
2454 return false;
2455 }
2456 RegisterTypesDetected[static_cast<int>(otherRegType)] = true;
2457 }
2458 }
2459 return true;
2460}
2461
2463 Decl *D, RegisterType RegType,
2464 bool SpecifiedSpace) {
2465
2466 // exactly one of these two types should be set
2467 assert(((isa<VarDecl>(D) && !isa<HLSLBufferDecl>(D)) ||
2468 (!isa<VarDecl>(D) && isa<HLSLBufferDecl>(D))) &&
2469 "expecting VarDecl or HLSLBufferDecl");
2470
2471 // check if the declaration contains resource matching the register type
2472 if (!DiagnoseLocalRegisterBinding(S, ArgLoc, D, RegType, SpecifiedSpace))
2473 return false;
2474
2475 // next, if multiple register annotations exist, check that none conflict.
2476 return ValidateMultipleRegisterAnnotations(S, D, RegType);
2477}
2478
2479// return false if the slot count exceeds the limit, true otherwise
2480static bool AccumulateHLSLResourceSlots(QualType Ty, uint64_t &StartSlot,
2481 const uint64_t &Limit,
2482 const ResourceClass ResClass,
2483 ASTContext &Ctx,
2484 uint64_t ArrayCount = 1) {
2485 Ty = Ty.getCanonicalType();
2486 const Type *T = Ty.getTypePtr();
2487
2488 // Early exit if already overflowed
2489 if (StartSlot > Limit)
2490 return false;
2491
2492 // Case 1: array type
2493 if (const auto *AT = dyn_cast<ArrayType>(T)) {
2494 uint64_t Count = 1;
2495
2496 if (const auto *CAT = dyn_cast<ConstantArrayType>(AT))
2497 Count = CAT->getSize().getZExtValue();
2498
2499 QualType ElemTy = AT->getElementType();
2500 return AccumulateHLSLResourceSlots(ElemTy, StartSlot, Limit, ResClass, Ctx,
2501 ArrayCount * Count);
2502 }
2503
2504 // Case 2: resource leaf
2505 if (auto ResTy = dyn_cast<HLSLAttributedResourceType>(T)) {
2506 // First ensure this resource counts towards the corresponding
2507 // register type limit.
2508 if (ResTy->getAttrs().ResourceClass != ResClass)
2509 return true;
2510
2511 // Validate highest slot used
2512 uint64_t EndSlot = StartSlot + ArrayCount - 1;
2513 if (EndSlot > Limit)
2514 return false;
2515
2516 // Advance SlotCount past the consumed range
2517 StartSlot = EndSlot + 1;
2518 return true;
2519 }
2520
2521 // Case 3: struct / record
2522 if (const auto *RT = dyn_cast<RecordType>(T)) {
2523 const RecordDecl *RD = RT->getDecl();
2524
2525 if (const auto *CXXRD = dyn_cast<CXXRecordDecl>(RD)) {
2526 for (const CXXBaseSpecifier &Base : CXXRD->bases()) {
2527 if (!AccumulateHLSLResourceSlots(Base.getType(), StartSlot, Limit,
2528 ResClass, Ctx, ArrayCount))
2529 return false;
2530 }
2531 }
2532
2533 for (const FieldDecl *Field : RD->fields()) {
2534 if (!AccumulateHLSLResourceSlots(Field->getType(), StartSlot, Limit,
2535 ResClass, Ctx, ArrayCount))
2536 return false;
2537 }
2538
2539 return true;
2540 }
2541
2542 // Case 4: everything else
2543 return true;
2544}
2545
2546// return true if there is something invalid, false otherwise
2547static bool ValidateRegisterNumber(uint64_t SlotNum, Decl *TheDecl,
2548 ASTContext &Ctx, RegisterType RegTy) {
2549 const uint64_t Limit = UINT32_MAX;
2550 if (SlotNum > Limit)
2551 return true;
2552
2553 // after verifying the number doesn't exceed uint32max, we don't need
2554 // to look further into c or i register types
2555 if (RegTy == RegisterType::C || RegTy == RegisterType::I)
2556 return false;
2557
2558 if (VarDecl *VD = dyn_cast<VarDecl>(TheDecl)) {
2559 uint64_t BaseSlot = SlotNum;
2560
2561 if (!AccumulateHLSLResourceSlots(VD->getType(), SlotNum, Limit,
2562 getResourceClass(RegTy), Ctx))
2563 return true;
2564
2565 // After AccumulateHLSLResourceSlots runs, SlotNum is now
2566 // the first free slot; last used was SlotNum - 1
2567 return (BaseSlot > Limit);
2568 }
2569 // handle the cbuffer/tbuffer case
2570 if (isa<HLSLBufferDecl>(TheDecl))
2571 // resources cannot be put within a cbuffer, so no need
2572 // to analyze the structure since the register number
2573 // won't be pushed any higher.
2574 return (SlotNum > Limit);
2575
2576 // we don't expect any other decl type, so fail
2577 llvm_unreachable("unexpected decl type");
2578}
2579
2581 if (VarDecl *VD = dyn_cast<VarDecl>(TheDecl)) {
2582 QualType Ty = VD->getType();
2583 if (const auto *IAT = dyn_cast<IncompleteArrayType>(Ty))
2584 Ty = IAT->getElementType();
2585 if (SemaRef.RequireCompleteType(TheDecl->getBeginLoc(), Ty,
2586 diag::err_incomplete_type))
2587 return;
2588 }
2589
2590 StringRef Slot = "";
2591 StringRef Space = "";
2592 SourceLocation SlotLoc, SpaceLoc;
2593
2594 if (!AL.isArgIdent(0)) {
2595 Diag(AL.getLoc(), diag::err_attribute_argument_type)
2596 << AL << AANT_ArgumentIdentifier;
2597 return;
2598 }
2599 IdentifierLoc *Loc = AL.getArgAsIdent(0);
2600
2601 if (AL.getNumArgs() == 2) {
2602 Slot = Loc->getIdentifierInfo()->getName();
2603 SlotLoc = Loc->getLoc();
2604 if (!AL.isArgIdent(1)) {
2605 Diag(AL.getLoc(), diag::err_attribute_argument_type)
2606 << AL << AANT_ArgumentIdentifier;
2607 return;
2608 }
2609 Loc = AL.getArgAsIdent(1);
2610 Space = Loc->getIdentifierInfo()->getName();
2611 SpaceLoc = Loc->getLoc();
2612 } else {
2613 StringRef Str = Loc->getIdentifierInfo()->getName();
2614 if (Str.starts_with("space")) {
2615 Space = Str;
2616 SpaceLoc = Loc->getLoc();
2617 } else {
2618 Slot = Str;
2619 SlotLoc = Loc->getLoc();
2620 Space = "space0";
2621 }
2622 }
2623
2624 RegisterType RegType = RegisterType::SRV;
2625 std::optional<unsigned> SlotNum;
2626 unsigned SpaceNum = 0;
2627
2628 // Validate slot
2629 if (!Slot.empty()) {
2630 if (!convertToRegisterType(Slot, &RegType)) {
2631 Diag(SlotLoc, diag::err_hlsl_binding_type_invalid) << Slot.substr(0, 1);
2632 return;
2633 }
2634 if (RegType == RegisterType::I) {
2635 Diag(SlotLoc, diag::warn_hlsl_deprecated_register_type_i);
2636 return;
2637 }
2638 const StringRef SlotNumStr = Slot.substr(1);
2639
2640 uint64_t N;
2641
2642 // validate that the slot number is a non-empty number
2643 if (SlotNumStr.getAsInteger(10, N)) {
2644 Diag(SlotLoc, diag::err_hlsl_unsupported_register_number);
2645 return;
2646 }
2647
2648 // Validate register number. It should not exceed UINT32_MAX,
2649 // including if the resource type is an array that starts
2650 // before UINT32_MAX, but ends afterwards.
2651 if (ValidateRegisterNumber(N, TheDecl, getASTContext(), RegType)) {
2652 Diag(SlotLoc, diag::err_hlsl_register_number_too_large);
2653 return;
2654 }
2655
2656 // the slot number has been validated and does not exceed UINT32_MAX
2657 SlotNum = (unsigned)N;
2658 }
2659
2660 // Validate space
2661 if (!Space.starts_with("space")) {
2662 Diag(SpaceLoc, diag::err_hlsl_expected_space) << Space;
2663 return;
2664 }
2665 StringRef SpaceNumStr = Space.substr(5);
2666 if (SpaceNumStr.getAsInteger(10, SpaceNum)) {
2667 Diag(SpaceLoc, diag::err_hlsl_expected_space) << Space;
2668 return;
2669 }
2670
2671 // If we have slot, diagnose it is the right register type for the decl
2672 if (SlotNum.has_value())
2673 if (!DiagnoseHLSLRegisterAttribute(SemaRef, SlotLoc, TheDecl, RegType,
2674 !SpaceLoc.isInvalid()))
2675 return;
2676
2677 HLSLResourceBindingAttr *NewAttr =
2678 HLSLResourceBindingAttr::Create(getASTContext(), Slot, Space, AL);
2679 if (NewAttr) {
2680 NewAttr->setBinding(RegType, SlotNum, SpaceNum);
2681 TheDecl->addAttr(NewAttr);
2682 }
2683}
2684
2686 HLSLParamModifierAttr *NewAttr = mergeParamModifierAttr(
2687 D, AL,
2688 static_cast<HLSLParamModifierAttr::Spelling>(AL.getSemanticSpelling()));
2689 if (NewAttr)
2690 D->addAttr(NewAttr);
2691}
2692
2693namespace {
2694
2695/// This class implements HLSL availability diagnostics for default
2696/// and relaxed mode
2697///
2698/// The goal of this diagnostic is to emit an error or warning when an
2699/// unavailable API is found in code that is reachable from the shader
2700/// entry function or from an exported function (when compiling a shader
2701/// library).
2702///
2703/// This is done by traversing the AST of all shader entry point functions
2704/// and of all exported functions, and any functions that are referenced
2705/// from this AST. In other words, any functions that are reachable from
2706/// the entry points.
2707class DiagnoseHLSLAvailability : public DynamicRecursiveASTVisitor {
2708 Sema &SemaRef;
2709
2710 // Stack of functions to be scaned
2712
2713 // Tracks which environments functions have been scanned in.
2714 //
2715 // Maps FunctionDecl to an unsigned number that represents the set of shader
2716 // environments the function has been scanned for.
2717 // The llvm::Triple::EnvironmentType enum values for shader stages guaranteed
2718 // to be numbered from llvm::Triple::Pixel to llvm::Triple::Amplification
2719 // (verified by static_asserts in Triple.cpp), we can use it to index
2720 // individual bits in the set, as long as we shift the values to start with 0
2721 // by subtracting the value of llvm::Triple::Pixel first.
2722 //
2723 // The N'th bit in the set will be set if the function has been scanned
2724 // in shader environment whose llvm::Triple::EnvironmentType integer value
2725 // equals (llvm::Triple::Pixel + N).
2726 //
2727 // For example, if a function has been scanned in compute and pixel stage
2728 // environment, the value will be 0x21 (100001 binary) because:
2729 //
2730 // (int)(llvm::Triple::Pixel - llvm::Triple::Pixel) == 0
2731 // (int)(llvm::Triple::Compute - llvm::Triple::Pixel) == 5
2732 //
2733 // A FunctionDecl is mapped to 0 (or not included in the map) if it has not
2734 // been scanned in any environment.
2735 llvm::DenseMap<const FunctionDecl *, unsigned> ScannedDecls;
2736
2737 // Do not access these directly, use the get/set methods below to make
2738 // sure the values are in sync
2739 llvm::Triple::EnvironmentType CurrentShaderEnvironment;
2740 unsigned CurrentShaderStageBit;
2741
2742 // True if scanning a function that was already scanned in a different
2743 // shader stage context, and therefore we should not report issues that
2744 // depend only on shader model version because they would be duplicate.
2745 bool ReportOnlyShaderStageIssues;
2746
2747 // Helper methods for dealing with current stage context / environment
2748 void SetShaderStageContext(llvm::Triple::EnvironmentType ShaderType) {
2749 static_assert(sizeof(unsigned) >= 4);
2750 assert(HLSLShaderAttr::isValidShaderType(ShaderType));
2751 assert((unsigned)(ShaderType - llvm::Triple::Pixel) < 31 &&
2752 "ShaderType is too big for this bitmap"); // 31 is reserved for
2753 // "unknown"
2754
2755 unsigned bitmapIndex = ShaderType - llvm::Triple::Pixel;
2756 CurrentShaderEnvironment = ShaderType;
2757 CurrentShaderStageBit = (1 << bitmapIndex);
2758 }
2759
2760 void SetUnknownShaderStageContext() {
2761 CurrentShaderEnvironment = llvm::Triple::UnknownEnvironment;
2762 CurrentShaderStageBit = (1 << 31);
2763 }
2764
2765 llvm::Triple::EnvironmentType GetCurrentShaderEnvironment() const {
2766 return CurrentShaderEnvironment;
2767 }
2768
2769 bool InUnknownShaderStageContext() const {
2770 return CurrentShaderEnvironment == llvm::Triple::UnknownEnvironment;
2771 }
2772
2773 // Helper methods for dealing with shader stage bitmap
2774 void AddToScannedFunctions(const FunctionDecl *FD) {
2775 unsigned &ScannedStages = ScannedDecls[FD];
2776 ScannedStages |= CurrentShaderStageBit;
2777 }
2778
2779 unsigned GetScannedStages(const FunctionDecl *FD) { return ScannedDecls[FD]; }
2780
2781 bool WasAlreadyScannedInCurrentStage(const FunctionDecl *FD) {
2782 return WasAlreadyScannedInCurrentStage(GetScannedStages(FD));
2783 }
2784
2785 bool WasAlreadyScannedInCurrentStage(unsigned ScannerStages) {
2786 return ScannerStages & CurrentShaderStageBit;
2787 }
2788
2789 static bool NeverBeenScanned(unsigned ScannedStages) {
2790 return ScannedStages == 0;
2791 }
2792
2793 // Scanning methods
2794 void HandleFunctionOrMethodRef(FunctionDecl *FD, Expr *RefExpr);
2795 void CheckDeclAvailability(NamedDecl *D, const AvailabilityAttr *AA,
2796 SourceRange Range);
2797 const AvailabilityAttr *FindAvailabilityAttr(const Decl *D);
2798 bool HasMatchingEnvironmentOrNone(const AvailabilityAttr *AA);
2799
2800public:
2801 DiagnoseHLSLAvailability(Sema &SemaRef)
2802 : SemaRef(SemaRef),
2803 CurrentShaderEnvironment(llvm::Triple::UnknownEnvironment),
2804 CurrentShaderStageBit(0), ReportOnlyShaderStageIssues(false) {}
2805
2806 // AST traversal methods
2807 void RunOnTranslationUnit(const TranslationUnitDecl *TU);
2808 void RunOnFunction(const FunctionDecl *FD);
2809
2810 bool VisitDeclRefExpr(DeclRefExpr *DRE) override {
2811 FunctionDecl *FD = llvm::dyn_cast<FunctionDecl>(DRE->getDecl());
2812 if (FD)
2813 HandleFunctionOrMethodRef(FD, DRE);
2814 return true;
2815 }
2816
2817 bool VisitMemberExpr(MemberExpr *ME) override {
2818 FunctionDecl *FD = llvm::dyn_cast<FunctionDecl>(ME->getMemberDecl());
2819 if (FD)
2820 HandleFunctionOrMethodRef(FD, ME);
2821 return true;
2822 }
2823};
2824
2825void DiagnoseHLSLAvailability::HandleFunctionOrMethodRef(FunctionDecl *FD,
2826 Expr *RefExpr) {
2827 assert((isa<DeclRefExpr>(RefExpr) || isa<MemberExpr>(RefExpr)) &&
2828 "expected DeclRefExpr or MemberExpr");
2829
2830 // has a definition -> add to stack to be scanned
2831 const FunctionDecl *FDWithBody = nullptr;
2832 if (FD->hasBody(FDWithBody)) {
2833 if (!WasAlreadyScannedInCurrentStage(FDWithBody))
2834 DeclsToScan.push_back(FDWithBody);
2835 return;
2836 }
2837
2838 // no body -> diagnose availability
2839 const AvailabilityAttr *AA = FindAvailabilityAttr(FD);
2840 if (AA)
2841 CheckDeclAvailability(
2842 FD, AA, SourceRange(RefExpr->getBeginLoc(), RefExpr->getEndLoc()));
2843}
2844
2845void DiagnoseHLSLAvailability::RunOnTranslationUnit(
2846 const TranslationUnitDecl *TU) {
2847
2848 // Iterate over all shader entry functions and library exports, and for those
2849 // that have a body (definiton), run diag scan on each, setting appropriate
2850 // shader environment context based on whether it is a shader entry function
2851 // or an exported function. Exported functions can be in namespaces and in
2852 // export declarations so we need to scan those declaration contexts as well.
2854 DeclContextsToScan.push_back(TU);
2855
2856 while (!DeclContextsToScan.empty()) {
2857 const DeclContext *DC = DeclContextsToScan.pop_back_val();
2858 for (auto &D : DC->decls()) {
2859 // do not scan implicit declaration generated by the implementation
2860 if (D->isImplicit())
2861 continue;
2862
2863 // for namespace or export declaration add the context to the list to be
2864 // scanned later
2865 if (llvm::dyn_cast<NamespaceDecl>(D) || llvm::dyn_cast<ExportDecl>(D)) {
2866 DeclContextsToScan.push_back(llvm::dyn_cast<DeclContext>(D));
2867 continue;
2868 }
2869
2870 // skip over other decls or function decls without body
2871 const FunctionDecl *FD = llvm::dyn_cast<FunctionDecl>(D);
2872 if (!FD || !FD->isThisDeclarationADefinition())
2873 continue;
2874
2875 // shader entry point
2876 if (HLSLShaderAttr *ShaderAttr = FD->getAttr<HLSLShaderAttr>()) {
2877 SetShaderStageContext(ShaderAttr->getType());
2878 RunOnFunction(FD);
2879 continue;
2880 }
2881 // exported library function
2882 // FIXME: replace this loop with external linkage check once issue #92071
2883 // is resolved
2884 bool isExport = FD->isInExportDeclContext();
2885 if (!isExport) {
2886 for (const auto *Redecl : FD->redecls()) {
2887 if (Redecl->isInExportDeclContext()) {
2888 isExport = true;
2889 break;
2890 }
2891 }
2892 }
2893 if (isExport) {
2894 SetUnknownShaderStageContext();
2895 RunOnFunction(FD);
2896 continue;
2897 }
2898 }
2899 }
2900}
2901
2902void DiagnoseHLSLAvailability::RunOnFunction(const FunctionDecl *FD) {
2903 assert(DeclsToScan.empty() && "DeclsToScan should be empty");
2904 DeclsToScan.push_back(FD);
2905
2906 while (!DeclsToScan.empty()) {
2907 // Take one decl from the stack and check it by traversing its AST.
2908 // For any CallExpr found during the traversal add it's callee to the top of
2909 // the stack to be processed next. Functions already processed are stored in
2910 // ScannedDecls.
2911 const FunctionDecl *FD = DeclsToScan.pop_back_val();
2912
2913 // Decl was already scanned
2914 const unsigned ScannedStages = GetScannedStages(FD);
2915 if (WasAlreadyScannedInCurrentStage(ScannedStages))
2916 continue;
2917
2918 ReportOnlyShaderStageIssues = !NeverBeenScanned(ScannedStages);
2919
2920 AddToScannedFunctions(FD);
2921 TraverseStmt(FD->getBody());
2922 }
2923}
2924
2925bool DiagnoseHLSLAvailability::HasMatchingEnvironmentOrNone(
2926 const AvailabilityAttr *AA) {
2927 const IdentifierInfo *IIEnvironment = AA->getEnvironment();
2928 if (!IIEnvironment)
2929 return true;
2930
2931 llvm::Triple::EnvironmentType CurrentEnv = GetCurrentShaderEnvironment();
2932 if (CurrentEnv == llvm::Triple::UnknownEnvironment)
2933 return false;
2934
2935 llvm::Triple::EnvironmentType AttrEnv =
2936 AvailabilityAttr::getEnvironmentType(IIEnvironment->getName());
2937
2938 return CurrentEnv == AttrEnv;
2939}
2940
2941const AvailabilityAttr *
2942DiagnoseHLSLAvailability::FindAvailabilityAttr(const Decl *D) {
2943 AvailabilityAttr const *PartialMatch = nullptr;
2944 // Check each AvailabilityAttr to find the one for this platform.
2945 // For multiple attributes with the same platform try to find one for this
2946 // environment.
2947 for (const auto *A : D->attrs()) {
2948 if (const auto *Avail = dyn_cast<AvailabilityAttr>(A)) {
2949 StringRef AttrPlatform = Avail->getPlatform()->getName();
2950 StringRef TargetPlatform =
2952
2953 // Match the platform name.
2954 if (AttrPlatform == TargetPlatform) {
2955 // Find the best matching attribute for this environment
2956 if (HasMatchingEnvironmentOrNone(Avail))
2957 return Avail;
2958 PartialMatch = Avail;
2959 }
2960 }
2961 }
2962 return PartialMatch;
2963}
2964
2965// Check availability against target shader model version and current shader
2966// stage and emit diagnostic
2967void DiagnoseHLSLAvailability::CheckDeclAvailability(NamedDecl *D,
2968 const AvailabilityAttr *AA,
2969 SourceRange Range) {
2970
2971 const IdentifierInfo *IIEnv = AA->getEnvironment();
2972
2973 if (!IIEnv) {
2974 // The availability attribute does not have environment -> it depends only
2975 // on shader model version and not on specific the shader stage.
2976
2977 // Skip emitting the diagnostics if the diagnostic mode is set to
2978 // strict (-fhlsl-strict-availability) because all relevant diagnostics
2979 // were already emitted in the DiagnoseUnguardedAvailability scan
2980 // (SemaAvailability.cpp).
2981 if (SemaRef.getLangOpts().HLSLStrictAvailability)
2982 return;
2983
2984 // Do not report shader-stage-independent issues if scanning a function
2985 // that was already scanned in a different shader stage context (they would
2986 // be duplicate)
2987 if (ReportOnlyShaderStageIssues)
2988 return;
2989
2990 } else {
2991 // The availability attribute has environment -> we need to know
2992 // the current stage context to property diagnose it.
2993 if (InUnknownShaderStageContext())
2994 return;
2995 }
2996
2997 // Check introduced version and if environment matches
2998 bool EnvironmentMatches = HasMatchingEnvironmentOrNone(AA);
2999 VersionTuple Introduced = AA->getIntroduced();
3000 VersionTuple TargetVersion =
3002
3003 if (TargetVersion >= Introduced && EnvironmentMatches)
3004 return;
3005
3006 // Emit diagnostic message
3007 const TargetInfo &TI = SemaRef.getASTContext().getTargetInfo();
3008 llvm::StringRef PlatformName(
3009 AvailabilityAttr::getPrettyPlatformName(TI.getPlatformName()));
3010
3011 llvm::StringRef CurrentEnvStr =
3012 llvm::Triple::getEnvironmentTypeName(GetCurrentShaderEnvironment());
3013
3014 llvm::StringRef AttrEnvStr =
3015 AA->getEnvironment() ? AA->getEnvironment()->getName() : "";
3016 bool UseEnvironment = !AttrEnvStr.empty();
3017
3018 if (EnvironmentMatches) {
3019 SemaRef.Diag(Range.getBegin(), diag::warn_hlsl_availability)
3020 << Range << D << PlatformName << Introduced.getAsString()
3021 << UseEnvironment << CurrentEnvStr;
3022 } else {
3023 SemaRef.Diag(Range.getBegin(), diag::warn_hlsl_availability_unavailable)
3024 << Range << D;
3025 }
3026
3027 SemaRef.Diag(D->getLocation(), diag::note_partial_availability_specified_here)
3028 << D << PlatformName << Introduced.getAsString()
3029 << SemaRef.Context.getTargetInfo().getPlatformMinVersion().getAsString()
3030 << UseEnvironment << AttrEnvStr << CurrentEnvStr;
3031}
3032
3033} // namespace
3034
3036 // process default CBuffer - create buffer layout struct and invoke codegenCGH
3037 if (!DefaultCBufferDecls.empty()) {
3039 SemaRef.getASTContext(), SemaRef.getCurLexicalContext(),
3040 DefaultCBufferDecls);
3041 addImplicitBindingAttrToDecl(SemaRef, DefaultCBuffer, RegisterType::CBuffer,
3043 SemaRef.getCurLexicalContext()->addDecl(DefaultCBuffer);
3045
3046 // Set HasValidPackoffset if any of the decls has a register(c#) annotation;
3047 for (const Decl *VD : DefaultCBufferDecls) {
3048 const HLSLResourceBindingAttr *RBA =
3049 VD->getAttr<HLSLResourceBindingAttr>();
3050 if (RBA && RBA->hasRegisterSlot() &&
3051 RBA->getRegisterType() == HLSLResourceBindingAttr::RegisterType::C) {
3052 DefaultCBuffer->setHasValidPackoffset(true);
3053 break;
3054 }
3055 }
3056
3057 DeclGroupRef DG(DefaultCBuffer);
3058 SemaRef.Consumer.HandleTopLevelDecl(DG);
3059 }
3060 diagnoseAvailabilityViolations(TU);
3061}
3062
3063void SemaHLSL::diagnoseAvailabilityViolations(TranslationUnitDecl *TU) {
3064 // Skip running the diagnostics scan if the diagnostic mode is
3065 // strict (-fhlsl-strict-availability) and the target shader stage is known
3066 // because all relevant diagnostics were already emitted in the
3067 // DiagnoseUnguardedAvailability scan (SemaAvailability.cpp).
3069 if (SemaRef.getLangOpts().HLSLStrictAvailability &&
3070 TI.getTriple().getEnvironment() != llvm::Triple::EnvironmentType::Library)
3071 return;
3072
3073 DiagnoseHLSLAvailability(SemaRef).RunOnTranslationUnit(TU);
3074}
3075
3076static bool CheckAllArgsHaveSameType(Sema *S, CallExpr *TheCall) {
3077 assert(TheCall->getNumArgs() > 1);
3078 QualType ArgTy0 = TheCall->getArg(0)->getType();
3079
3080 for (unsigned I = 1, N = TheCall->getNumArgs(); I < N; ++I) {
3082 ArgTy0, TheCall->getArg(I)->getType())) {
3083 S->Diag(TheCall->getBeginLoc(), diag::err_vec_builtin_incompatible_vector)
3084 << TheCall->getDirectCallee() << /*useAllTerminology*/ true
3085 << SourceRange(TheCall->getArg(0)->getBeginLoc(),
3086 TheCall->getArg(N - 1)->getEndLoc());
3087 return true;
3088 }
3089 }
3090 return false;
3091}
3092
3094 QualType ArgType = Arg->getType();
3096 S->Diag(Arg->getBeginLoc(), diag::err_typecheck_convert_incompatible)
3097 << ArgType << ExpectedType << 1 << 0 << 0;
3098 return true;
3099 }
3100 return false;
3101}
3102
3104 Sema *S, CallExpr *TheCall,
3105 llvm::function_ref<bool(Sema *S, SourceLocation Loc, int ArgOrdinal,
3106 clang::QualType PassedType)>
3107 Check) {
3108 for (unsigned I = 0; I < TheCall->getNumArgs(); ++I) {
3109 Expr *Arg = TheCall->getArg(I);
3110 if (Check(S, Arg->getBeginLoc(), I + 1, Arg->getType()))
3111 return true;
3112 }
3113 return false;
3114}
3115
3117 int ArgOrdinal,
3118 clang::QualType PassedType) {
3119 clang::QualType BaseType =
3120 PassedType->isVectorType()
3121 ? PassedType->castAs<clang::VectorType>()->getElementType()
3122 : PassedType;
3123 if (!BaseType->isFloat32Type())
3124 return S->Diag(Loc, diag::err_builtin_invalid_arg_type)
3125 << ArgOrdinal << /* scalar or vector of */ 5 << /* no int */ 0
3126 << /* float */ 1 << PassedType;
3127 return false;
3128}
3129
3131 int ArgOrdinal,
3132 clang::QualType PassedType) {
3133 clang::QualType BaseType =
3134 PassedType->isVectorType()
3135 ? PassedType->castAs<clang::VectorType>()->getElementType()
3136 : PassedType;
3137 if (!BaseType->isHalfType() && !BaseType->isFloat32Type())
3138 return S->Diag(Loc, diag::err_builtin_invalid_arg_type)
3139 << ArgOrdinal << /* scalar or vector of */ 5 << /* no int */ 0
3140 << /* half or float */ 2 << PassedType;
3141 return false;
3142}
3143
3144static bool CheckModifiableLValue(Sema *S, CallExpr *TheCall,
3145 unsigned ArgIndex) {
3146 auto *Arg = TheCall->getArg(ArgIndex);
3147 SourceLocation OrigLoc = Arg->getExprLoc();
3148 if (Arg->IgnoreCasts()->isModifiableLvalue(S->Context, &OrigLoc) ==
3150 return false;
3151 S->Diag(OrigLoc, diag::error_hlsl_inout_lvalue) << Arg << 0;
3152 return true;
3153}
3154
3155static bool CheckNoDoubleVectors(Sema *S, SourceLocation Loc, int ArgOrdinal,
3156 clang::QualType PassedType) {
3157 const auto *VecTy = PassedType->getAs<VectorType>();
3158 if (!VecTy)
3159 return false;
3160
3161 if (VecTy->getElementType()->isDoubleType())
3162 return S->Diag(Loc, diag::err_builtin_invalid_arg_type)
3163 << ArgOrdinal << /* scalar */ 1 << /* no int */ 0 << /* fp */ 1
3164 << PassedType;
3165 return false;
3166}
3167
3169 int ArgOrdinal,
3170 clang::QualType PassedType) {
3171 if (!PassedType->hasIntegerRepresentation() &&
3172 !PassedType->hasFloatingRepresentation())
3173 return S->Diag(Loc, diag::err_builtin_invalid_arg_type)
3174 << ArgOrdinal << /* scalar or vector of */ 5 << /* integer */ 1
3175 << /* fp */ 1 << PassedType;
3176 return false;
3177}
3178
3180 int ArgOrdinal,
3181 clang::QualType PassedType) {
3182 if (auto *VecTy = PassedType->getAs<VectorType>())
3183 if (VecTy->getElementType()->isUnsignedIntegerType())
3184 return false;
3185
3186 return S->Diag(Loc, diag::err_builtin_invalid_arg_type)
3187 << ArgOrdinal << /* vector of */ 4 << /* uint */ 3 << /* no fp */ 0
3188 << PassedType;
3189}
3190
3191// checks for unsigned ints of all sizes
3193 int ArgOrdinal,
3194 clang::QualType PassedType) {
3195 if (!PassedType->hasUnsignedIntegerRepresentation())
3196 return S->Diag(Loc, diag::err_builtin_invalid_arg_type)
3197 << ArgOrdinal << /* scalar or vector of */ 5 << /* unsigned int */ 3
3198 << /* no fp */ 0 << PassedType;
3199 return false;
3200}
3201
3202static bool CheckExpectedBitWidth(Sema *S, CallExpr *TheCall,
3203 unsigned ArgOrdinal, unsigned Width) {
3204 QualType ArgTy = TheCall->getArg(0)->getType();
3205 if (auto *VTy = ArgTy->getAs<VectorType>())
3206 ArgTy = VTy->getElementType();
3207 // ensure arg type has expected bit width
3208 uint64_t ElementBitCount =
3210 if (ElementBitCount != Width) {
3211 S->Diag(TheCall->getArg(0)->getBeginLoc(),
3212 diag::err_integer_incorrect_bit_count)
3213 << Width << ElementBitCount;
3214 return true;
3215 }
3216 return false;
3217}
3218
3220 QualType ReturnType) {
3221 auto *VecTyA = TheCall->getArg(0)->getType()->getAs<VectorType>();
3222 if (VecTyA)
3223 ReturnType =
3224 S->Context.getExtVectorType(ReturnType, VecTyA->getNumElements());
3225
3226 TheCall->setType(ReturnType);
3227}
3228
3229static bool CheckScalarOrVector(Sema *S, CallExpr *TheCall, QualType Scalar,
3230 unsigned ArgIndex) {
3231 assert(TheCall->getNumArgs() >= ArgIndex);
3232 QualType ArgType = TheCall->getArg(ArgIndex)->getType();
3233 auto *VTy = ArgType->getAs<VectorType>();
3234 // not the scalar or vector<scalar>
3235 if (!(S->Context.hasSameUnqualifiedType(ArgType, Scalar) ||
3236 (VTy &&
3237 S->Context.hasSameUnqualifiedType(VTy->getElementType(), Scalar)))) {
3238 S->Diag(TheCall->getArg(0)->getBeginLoc(),
3239 diag::err_typecheck_expect_scalar_or_vector)
3240 << ArgType << Scalar;
3241 return true;
3242 }
3243 return false;
3244}
3245
3247 QualType Scalar, unsigned ArgIndex) {
3248 assert(TheCall->getNumArgs() > ArgIndex);
3249
3250 Expr *Arg = TheCall->getArg(ArgIndex);
3251 QualType ArgType = Arg->getType();
3252
3253 // Scalar: T
3254 if (S->Context.hasSameUnqualifiedType(ArgType, Scalar))
3255 return false;
3256
3257 // Vector: vector<T>
3258 if (const auto *VTy = ArgType->getAs<VectorType>()) {
3259 if (S->Context.hasSameUnqualifiedType(VTy->getElementType(), Scalar))
3260 return false;
3261 }
3262
3263 // Matrix: ConstantMatrixType with element type T
3264 if (const auto *MTy = ArgType->getAs<ConstantMatrixType>()) {
3265 if (S->Context.hasSameUnqualifiedType(MTy->getElementType(), Scalar))
3266 return false;
3267 }
3268
3269 // Not a scalar/vector/matrix-of-scalar
3270 S->Diag(Arg->getBeginLoc(),
3271 diag::err_typecheck_expect_scalar_or_vector_or_matrix)
3272 << ArgType << Scalar;
3273 return true;
3274}
3275
3276static bool CheckAnyScalarOrVector(Sema *S, CallExpr *TheCall,
3277 unsigned ArgIndex) {
3278 assert(TheCall->getNumArgs() >= ArgIndex);
3279 QualType ArgType = TheCall->getArg(ArgIndex)->getType();
3280 auto *VTy = ArgType->getAs<VectorType>();
3281 // not the scalar or vector<scalar>
3282 if (!(ArgType->isScalarType() ||
3283 (VTy && VTy->getElementType()->isScalarType()))) {
3284 S->Diag(TheCall->getArg(0)->getBeginLoc(),
3285 diag::err_typecheck_expect_any_scalar_or_vector)
3286 << ArgType << 1;
3287 return true;
3288 }
3289 return false;
3290}
3291
3292// Check that the argument is not a bool or vector<bool>
3293// Returns true on error
3295 unsigned ArgIndex) {
3296 QualType BoolType = S->getASTContext().BoolTy;
3297 assert(ArgIndex < TheCall->getNumArgs());
3298 QualType ArgType = TheCall->getArg(ArgIndex)->getType();
3299 auto *VTy = ArgType->getAs<VectorType>();
3300 // is the bool or vector<bool>
3301 if (S->Context.hasSameUnqualifiedType(ArgType, BoolType) ||
3302 (VTy &&
3303 S->Context.hasSameUnqualifiedType(VTy->getElementType(), BoolType))) {
3304 S->Diag(TheCall->getArg(0)->getBeginLoc(),
3305 diag::err_typecheck_expect_any_scalar_or_vector)
3306 << ArgType << 0;
3307 return true;
3308 }
3309 return false;
3310}
3311
3312static bool CheckWaveActive(Sema *S, CallExpr *TheCall) {
3313 if (CheckNotBoolScalarOrVector(S, TheCall, 0))
3314 return true;
3315 return false;
3316}
3317
3318static bool CheckWavePrefix(Sema *S, CallExpr *TheCall) {
3319 if (CheckNotBoolScalarOrVector(S, TheCall, 0))
3320 return true;
3321 return false;
3322}
3323
3324static bool CheckBoolSelect(Sema *S, CallExpr *TheCall) {
3325 assert(TheCall->getNumArgs() == 3);
3326 Expr *Arg1 = TheCall->getArg(1);
3327 Expr *Arg2 = TheCall->getArg(2);
3328 if (!S->Context.hasSameUnqualifiedType(Arg1->getType(), Arg2->getType())) {
3329 S->Diag(TheCall->getBeginLoc(),
3330 diag::err_typecheck_call_different_arg_types)
3331 << Arg1->getType() << Arg2->getType() << Arg1->getSourceRange()
3332 << Arg2->getSourceRange();
3333 return true;
3334 }
3335
3336 TheCall->setType(Arg1->getType());
3337 return false;
3338}
3339
3340static bool CheckVectorSelect(Sema *S, CallExpr *TheCall) {
3341 assert(TheCall->getNumArgs() == 3);
3342 Expr *Arg1 = TheCall->getArg(1);
3343 QualType Arg1Ty = Arg1->getType();
3344 Expr *Arg2 = TheCall->getArg(2);
3345 QualType Arg2Ty = Arg2->getType();
3346
3347 QualType Arg1ScalarTy = Arg1Ty;
3348 if (auto VTy = Arg1ScalarTy->getAs<VectorType>())
3349 Arg1ScalarTy = VTy->getElementType();
3350
3351 QualType Arg2ScalarTy = Arg2Ty;
3352 if (auto VTy = Arg2ScalarTy->getAs<VectorType>())
3353 Arg2ScalarTy = VTy->getElementType();
3354
3355 if (!S->Context.hasSameUnqualifiedType(Arg1ScalarTy, Arg2ScalarTy))
3356 S->Diag(Arg1->getBeginLoc(), diag::err_hlsl_builtin_scalar_vector_mismatch)
3357 << /* second and third */ 1 << TheCall->getCallee() << Arg1Ty << Arg2Ty;
3358
3359 QualType Arg0Ty = TheCall->getArg(0)->getType();
3360 unsigned Arg0Length = Arg0Ty->getAs<VectorType>()->getNumElements();
3361 unsigned Arg1Length = Arg1Ty->isVectorType()
3362 ? Arg1Ty->getAs<VectorType>()->getNumElements()
3363 : 0;
3364 unsigned Arg2Length = Arg2Ty->isVectorType()
3365 ? Arg2Ty->getAs<VectorType>()->getNumElements()
3366 : 0;
3367 if (Arg1Length > 0 && Arg0Length != Arg1Length) {
3368 S->Diag(TheCall->getBeginLoc(),
3369 diag::err_typecheck_vector_lengths_not_equal)
3370 << Arg0Ty << Arg1Ty << TheCall->getArg(0)->getSourceRange()
3371 << Arg1->getSourceRange();
3372 return true;
3373 }
3374
3375 if (Arg2Length > 0 && Arg0Length != Arg2Length) {
3376 S->Diag(TheCall->getBeginLoc(),
3377 diag::err_typecheck_vector_lengths_not_equal)
3378 << Arg0Ty << Arg2Ty << TheCall->getArg(0)->getSourceRange()
3379 << Arg2->getSourceRange();
3380 return true;
3381 }
3382
3383 TheCall->setType(
3384 S->getASTContext().getExtVectorType(Arg1ScalarTy, Arg0Length));
3385 return false;
3386}
3387
3389 Sema *S, CallExpr *TheCall, unsigned ArgIndex,
3390 llvm::function_ref<bool(const HLSLAttributedResourceType *ResType)> Check =
3391 nullptr) {
3392 assert(TheCall->getNumArgs() >= ArgIndex);
3393 QualType ArgType = TheCall->getArg(ArgIndex)->getType();
3394 const HLSLAttributedResourceType *ResTy =
3395 ArgType.getTypePtr()->getAs<HLSLAttributedResourceType>();
3396 if (!ResTy) {
3397 S->Diag(TheCall->getArg(ArgIndex)->getBeginLoc(),
3398 diag::err_typecheck_expect_hlsl_resource)
3399 << ArgType;
3400 return true;
3401 }
3402 if (Check && Check(ResTy)) {
3403 S->Diag(TheCall->getArg(ArgIndex)->getExprLoc(),
3404 diag::err_invalid_hlsl_resource_type)
3405 << ArgType;
3406 return true;
3407 }
3408 return false;
3409}
3410
3411static bool CheckVectorElementCount(Sema *S, QualType PassedType,
3412 QualType BaseType, unsigned ExpectedCount,
3413 SourceLocation Loc) {
3414 unsigned PassedCount = 1;
3415 if (const auto *VecTy = PassedType->getAs<VectorType>())
3416 PassedCount = VecTy->getNumElements();
3417
3418 if (PassedCount != ExpectedCount) {
3420 S->Context.getExtVectorType(BaseType, ExpectedCount);
3421 S->Diag(Loc, diag::err_typecheck_convert_incompatible)
3422 << PassedType << ExpectedType << 1 << 0 << 0;
3423 return true;
3424 }
3425 return false;
3426}
3427
3428enum class SampleKind { Sample, Bias, Grad, Level, Cmp, CmpLevelZero };
3429
3431 // Check the texture handle.
3432 if (CheckResourceHandle(&S, TheCall, 0,
3433 [](const HLSLAttributedResourceType *ResType) {
3434 return ResType->getAttrs().ResourceDimension ==
3435 llvm::dxil::ResourceDimension::Unknown;
3436 }))
3437 return true;
3438
3439 // Check the sampler handle.
3440 if (CheckResourceHandle(&S, TheCall, 1,
3441 [](const HLSLAttributedResourceType *ResType) {
3442 return ResType->getAttrs().ResourceClass !=
3443 llvm::hlsl::ResourceClass::Sampler;
3444 }))
3445 return true;
3446
3447 auto *ResourceTy =
3448 TheCall->getArg(0)->getType()->castAs<HLSLAttributedResourceType>();
3449
3450 // Check the location.
3451 unsigned ExpectedDim =
3452 getResourceDimensions(ResourceTy->getAttrs().ResourceDimension);
3453 if (CheckVectorElementCount(&S, TheCall->getArg(2)->getType(),
3454 S.Context.FloatTy, ExpectedDim,
3455 TheCall->getBeginLoc()))
3456 return true;
3457
3458 return false;
3459}
3460
3461static bool CheckGatherBuiltin(Sema &S, CallExpr *TheCall, bool IsCmp) {
3462 if (S.checkArgCountRange(TheCall, IsCmp ? 5 : 4, IsCmp ? 6 : 5))
3463 return true;
3464
3465 if (CheckTextureSamplerAndLocation(S, TheCall))
3466 return true;
3467
3468 unsigned NextIdx = 3;
3469 if (IsCmp) {
3470 // Check the compare value.
3471 QualType CmpTy = TheCall->getArg(NextIdx)->getType();
3472 if (!CmpTy->isFloatingType() || CmpTy->isVectorType()) {
3473 S.Diag(TheCall->getArg(NextIdx)->getBeginLoc(),
3474 diag::err_typecheck_convert_incompatible)
3475 << CmpTy << S.Context.FloatTy << 1 << 0 << 0;
3476 return true;
3477 }
3478 NextIdx++;
3479 }
3480
3481 // Check the component operand.
3482 Expr *ComponentArg = TheCall->getArg(NextIdx);
3483 QualType ComponentTy = ComponentArg->getType();
3484 if (!ComponentTy->isIntegerType() || ComponentTy->isVectorType()) {
3485 S.Diag(ComponentArg->getBeginLoc(),
3486 diag::err_typecheck_convert_incompatible)
3487 << ComponentTy << S.Context.UnsignedIntTy << 1 << 0 << 0;
3488 return true;
3489 }
3490
3491 // GatherCmp operations on Vulkan target must use component 0 (Red).
3492 if (IsCmp && S.getASTContext().getTargetInfo().getTriple().isSPIRV()) {
3493 std::optional<llvm::APSInt> ComponentOpt =
3494 ComponentArg->getIntegerConstantExpr(S.getASTContext());
3495 if (ComponentOpt) {
3496 int64_t ComponentVal = ComponentOpt->getSExtValue();
3497 if (ComponentVal != 0) {
3498 // Issue an error if the component is not 0 (Red).
3499 // 0 -> Red, 1 -> Green, 2 -> Blue, 3 -> Alpha
3500 assert(ComponentVal >= 0 && ComponentVal <= 3 &&
3501 "The component is not in the expected range.");
3502 S.Diag(ComponentArg->getBeginLoc(),
3503 diag::err_hlsl_gathercmp_invalid_component)
3504 << ComponentVal;
3505 return true;
3506 }
3507 }
3508 }
3509
3510 NextIdx++;
3511
3512 // Check the offset operand.
3513 const HLSLAttributedResourceType *ResourceTy =
3514 TheCall->getArg(0)->getType()->castAs<HLSLAttributedResourceType>();
3515 if (TheCall->getNumArgs() > NextIdx) {
3516 unsigned ExpectedDim =
3517 getResourceDimensions(ResourceTy->getAttrs().ResourceDimension);
3518 if (CheckVectorElementCount(&S, TheCall->getArg(NextIdx)->getType(),
3519 S.Context.IntTy, ExpectedDim,
3520 TheCall->getArg(NextIdx)->getBeginLoc()))
3521 return true;
3522 NextIdx++;
3523 }
3524
3525 assert(ResourceTy->hasContainedType() &&
3526 "Expecting a contained type for resource with a dimension "
3527 "attribute.");
3528 QualType ReturnType = ResourceTy->getContainedType();
3529
3530 if (IsCmp) {
3531 if (!ReturnType->hasFloatingRepresentation()) {
3532 S.Diag(TheCall->getBeginLoc(), diag::err_hlsl_samplecmp_requires_float);
3533 return true;
3534 }
3535 }
3536
3537 if (const auto *VecTy = ReturnType->getAs<VectorType>())
3538 ReturnType = VecTy->getElementType();
3539 ReturnType = S.Context.getExtVectorType(ReturnType, 4);
3540
3541 TheCall->setType(ReturnType);
3542
3543 return false;
3544}
3545static bool CheckLoadLevelBuiltin(Sema &S, CallExpr *TheCall) {
3546 if (S.checkArgCountRange(TheCall, 2, 3))
3547 return true;
3548
3549 // Check the texture handle.
3550 if (CheckResourceHandle(&S, TheCall, 0,
3551 [](const HLSLAttributedResourceType *ResType) {
3552 return ResType->getAttrs().ResourceDimension ==
3553 llvm::dxil::ResourceDimension::Unknown;
3554 }))
3555 return true;
3556
3557 auto *ResourceTy =
3558 TheCall->getArg(0)->getType()->castAs<HLSLAttributedResourceType>();
3559
3560 // Check the location + lod (int3 for Texture2D).
3561 unsigned ExpectedDim =
3562 getResourceDimensions(ResourceTy->getAttrs().ResourceDimension);
3563 QualType CoordLODTy = TheCall->getArg(1)->getType();
3564 if (CheckVectorElementCount(&S, CoordLODTy, S.Context.IntTy, ExpectedDim + 1,
3565 TheCall->getArg(1)->getBeginLoc()))
3566 return true;
3567
3568 QualType EltTy = CoordLODTy;
3569 if (const auto *VTy = EltTy->getAs<VectorType>())
3570 EltTy = VTy->getElementType();
3571 if (!EltTy->isIntegerType()) {
3572 S.Diag(TheCall->getArg(1)->getBeginLoc(), diag::err_typecheck_expect_int)
3573 << CoordLODTy;
3574 return true;
3575 }
3576
3577 // Check the offset operand.
3578 if (TheCall->getNumArgs() > 2) {
3579 if (CheckVectorElementCount(&S, TheCall->getArg(2)->getType(),
3580 S.Context.IntTy, ExpectedDim,
3581 TheCall->getArg(2)->getBeginLoc()))
3582 return true;
3583 }
3584
3585 TheCall->setType(ResourceTy->getContainedType());
3586 return false;
3587}
3588
3589static bool CheckSamplingBuiltin(Sema &S, CallExpr *TheCall, SampleKind Kind) {
3590 unsigned MinArgs, MaxArgs;
3591 if (Kind == SampleKind::Sample) {
3592 MinArgs = 3;
3593 MaxArgs = 5;
3594 } else if (Kind == SampleKind::Bias) {
3595 MinArgs = 4;
3596 MaxArgs = 6;
3597 } else if (Kind == SampleKind::Grad) {
3598 MinArgs = 5;
3599 MaxArgs = 7;
3600 } else if (Kind == SampleKind::Level) {
3601 MinArgs = 4;
3602 MaxArgs = 5;
3603 } else if (Kind == SampleKind::Cmp) {
3604 MinArgs = 4;
3605 MaxArgs = 6;
3606 } else {
3607 assert(Kind == SampleKind::CmpLevelZero);
3608 MinArgs = 4;
3609 MaxArgs = 5;
3610 }
3611
3612 if (S.checkArgCountRange(TheCall, MinArgs, MaxArgs))
3613 return true;
3614
3615 if (CheckTextureSamplerAndLocation(S, TheCall))
3616 return true;
3617
3618 const HLSLAttributedResourceType *ResourceTy =
3619 TheCall->getArg(0)->getType()->castAs<HLSLAttributedResourceType>();
3620 unsigned ExpectedDim =
3621 getResourceDimensions(ResourceTy->getAttrs().ResourceDimension);
3622
3623 unsigned NextIdx = 3;
3624 if (Kind == SampleKind::Bias || Kind == SampleKind::Level ||
3625 Kind == SampleKind::Cmp || Kind == SampleKind::CmpLevelZero) {
3626 // Check the bias, lod level, or compare value, depending on the kind.
3627 // All of them must be a scalar float value.
3628 QualType BiasOrLODOrCmpTy = TheCall->getArg(NextIdx)->getType();
3629 if (!BiasOrLODOrCmpTy->isFloatingType() ||
3630 BiasOrLODOrCmpTy->isVectorType()) {
3631 S.Diag(TheCall->getArg(NextIdx)->getBeginLoc(),
3632 diag::err_typecheck_convert_incompatible)
3633 << BiasOrLODOrCmpTy << S.Context.FloatTy << 1 << 0 << 0;
3634 return true;
3635 }
3636 NextIdx++;
3637 } else if (Kind == SampleKind::Grad) {
3638 // Check the DDX operand.
3639 if (CheckVectorElementCount(&S, TheCall->getArg(NextIdx)->getType(),
3640 S.Context.FloatTy, ExpectedDim,
3641 TheCall->getArg(NextIdx)->getBeginLoc()))
3642 return true;
3643
3644 // Check the DDY operand.
3645 if (CheckVectorElementCount(&S, TheCall->getArg(NextIdx + 1)->getType(),
3646 S.Context.FloatTy, ExpectedDim,
3647 TheCall->getArg(NextIdx + 1)->getBeginLoc()))
3648 return true;
3649 NextIdx += 2;
3650 }
3651
3652 // Check the offset operand.
3653 if (TheCall->getNumArgs() > NextIdx) {
3654 if (CheckVectorElementCount(&S, TheCall->getArg(NextIdx)->getType(),
3655 S.Context.IntTy, ExpectedDim,
3656 TheCall->getArg(NextIdx)->getBeginLoc()))
3657 return true;
3658 NextIdx++;
3659 }
3660
3661 // Check the clamp operand.
3662 if (Kind != SampleKind::Level && Kind != SampleKind::CmpLevelZero &&
3663 TheCall->getNumArgs() > NextIdx) {
3664 QualType ClampTy = TheCall->getArg(NextIdx)->getType();
3665 if (!ClampTy->isFloatingType() || ClampTy->isVectorType()) {
3666 S.Diag(TheCall->getArg(NextIdx)->getBeginLoc(),
3667 diag::err_typecheck_convert_incompatible)
3668 << ClampTy << S.Context.FloatTy << 1 << 0 << 0;
3669 return true;
3670 }
3671 }
3672
3673 assert(ResourceTy->hasContainedType() &&
3674 "Expecting a contained type for resource with a dimension "
3675 "attribute.");
3676 QualType ReturnType = ResourceTy->getContainedType();
3677 if (Kind == SampleKind::Cmp || Kind == SampleKind::CmpLevelZero) {
3678 if (!ReturnType->hasFloatingRepresentation()) {
3679 S.Diag(TheCall->getBeginLoc(), diag::err_hlsl_samplecmp_requires_float);
3680 return true;
3681 }
3682 ReturnType = S.Context.FloatTy;
3683 }
3684 TheCall->setType(ReturnType);
3685
3686 return false;
3687}
3688
3689// Note: returning true in this case results in CheckBuiltinFunctionCall
3690// returning an ExprError
3691bool SemaHLSL::CheckBuiltinFunctionCall(unsigned BuiltinID, CallExpr *TheCall) {
3692 switch (BuiltinID) {
3693 case Builtin::BI__builtin_hlsl_adduint64: {
3694 if (SemaRef.checkArgCount(TheCall, 2))
3695 return true;
3696
3697 if (CheckAllArgTypesAreCorrect(&SemaRef, TheCall,
3699 return true;
3700
3701 // ensure arg integers are 32-bits
3702 if (CheckExpectedBitWidth(&SemaRef, TheCall, 0, 32))
3703 return true;
3704
3705 // ensure both args are vectors of total bit size of a multiple of 64
3706 auto *VTy = TheCall->getArg(0)->getType()->getAs<VectorType>();
3707 int NumElementsArg = VTy->getNumElements();
3708 if (NumElementsArg != 2 && NumElementsArg != 4) {
3709 SemaRef.Diag(TheCall->getBeginLoc(), diag::err_vector_incorrect_bit_count)
3710 << 1 /*a multiple of*/ << 64 << NumElementsArg * 32;
3711 return true;
3712 }
3713
3714 // ensure first arg and second arg have the same type
3715 if (CheckAllArgsHaveSameType(&SemaRef, TheCall))
3716 return true;
3717
3718 ExprResult A = TheCall->getArg(0);
3719 QualType ArgTyA = A.get()->getType();
3720 // return type is the same as the input type
3721 TheCall->setType(ArgTyA);
3722 break;
3723 }
3724 case Builtin::BI__builtin_hlsl_resource_getpointer: {
3725 if (SemaRef.checkArgCount(TheCall, 2) ||
3726 CheckResourceHandle(&SemaRef, TheCall, 0) ||
3727 CheckArgTypeMatches(&SemaRef, TheCall->getArg(1),
3728 SemaRef.getASTContext().UnsignedIntTy))
3729 return true;
3730
3731 auto *ResourceTy =
3732 TheCall->getArg(0)->getType()->castAs<HLSLAttributedResourceType>();
3733 QualType ContainedTy = ResourceTy->getContainedType();
3734 auto ReturnType =
3735 SemaRef.Context.getAddrSpaceQualType(ContainedTy, LangAS::hlsl_device);
3736 ReturnType = SemaRef.Context.getPointerType(ReturnType);
3737 TheCall->setType(ReturnType);
3738 TheCall->setValueKind(VK_LValue);
3739
3740 break;
3741 }
3742 case Builtin::BI__builtin_hlsl_resource_getpointer_typed: {
3743 if (SemaRef.checkArgCount(TheCall, 3) ||
3744 CheckResourceHandle(&SemaRef, TheCall, 0) ||
3745 CheckArgTypeMatches(&SemaRef, TheCall->getArg(1),
3746 SemaRef.getASTContext().UnsignedIntTy))
3747 return true;
3748
3749 QualType ElementTy = TheCall->getArg(2)->getType();
3750 assert(ElementTy->isPointerType() &&
3751 "expected pointer type for second argument");
3752 ElementTy = ElementTy->getPointeeType();
3753
3754 // Reject array types
3755 if (ElementTy->isArrayType())
3756 return SemaRef.Diag(
3757 cast<FunctionDecl>(SemaRef.CurContext)->getPointOfInstantiation(),
3758 diag::err_invalid_use_of_array_type);
3759
3760 auto ReturnType =
3761 SemaRef.Context.getAddrSpaceQualType(ElementTy, LangAS::hlsl_device);
3762 ReturnType = SemaRef.Context.getPointerType(ReturnType);
3763 TheCall->setType(ReturnType);
3764
3765 break;
3766 }
3767 case Builtin::BI__builtin_hlsl_resource_load_with_status: {
3768 if (SemaRef.checkArgCount(TheCall, 3) ||
3769 CheckResourceHandle(&SemaRef, TheCall, 0) ||
3770 CheckArgTypeMatches(&SemaRef, TheCall->getArg(1),
3771 SemaRef.getASTContext().UnsignedIntTy) ||
3772 CheckArgTypeMatches(&SemaRef, TheCall->getArg(2),
3773 SemaRef.getASTContext().UnsignedIntTy) ||
3774 CheckModifiableLValue(&SemaRef, TheCall, 2))
3775 return true;
3776
3777 auto *ResourceTy =
3778 TheCall->getArg(0)->getType()->castAs<HLSLAttributedResourceType>();
3779 QualType ReturnType = ResourceTy->getContainedType();
3780 TheCall->setType(ReturnType);
3781
3782 break;
3783 }
3784 case Builtin::BI__builtin_hlsl_resource_load_with_status_typed: {
3785 if (SemaRef.checkArgCount(TheCall, 4) ||
3786 CheckResourceHandle(&SemaRef, TheCall, 0) ||
3787 CheckArgTypeMatches(&SemaRef, TheCall->getArg(1),
3788 SemaRef.getASTContext().UnsignedIntTy) ||
3789 CheckArgTypeMatches(&SemaRef, TheCall->getArg(2),
3790 SemaRef.getASTContext().UnsignedIntTy) ||
3791 CheckModifiableLValue(&SemaRef, TheCall, 2))
3792 return true;
3793
3794 QualType ReturnType = TheCall->getArg(3)->getType();
3795 assert(ReturnType->isPointerType() &&
3796 "expected pointer type for second argument");
3797 ReturnType = ReturnType->getPointeeType();
3798
3799 // Reject array types
3800 if (ReturnType->isArrayType())
3801 return SemaRef.Diag(
3802 cast<FunctionDecl>(SemaRef.CurContext)->getPointOfInstantiation(),
3803 diag::err_invalid_use_of_array_type);
3804
3805 TheCall->setType(ReturnType);
3806
3807 break;
3808 }
3809 case Builtin::BI__builtin_hlsl_resource_load_level:
3810 return CheckLoadLevelBuiltin(SemaRef, TheCall);
3811 case Builtin::BI__builtin_hlsl_resource_sample:
3813 case Builtin::BI__builtin_hlsl_resource_sample_bias:
3815 case Builtin::BI__builtin_hlsl_resource_sample_grad:
3817 case Builtin::BI__builtin_hlsl_resource_sample_level:
3819 case Builtin::BI__builtin_hlsl_resource_sample_cmp:
3821 case Builtin::BI__builtin_hlsl_resource_sample_cmp_level_zero:
3823 case Builtin::BI__builtin_hlsl_resource_gather:
3824 return CheckGatherBuiltin(SemaRef, TheCall, /*IsCmp=*/false);
3825 case Builtin::BI__builtin_hlsl_resource_gather_cmp:
3826 return CheckGatherBuiltin(SemaRef, TheCall, /*IsCmp=*/true);
3827 case Builtin::BI__builtin_hlsl_resource_uninitializedhandle: {
3828 assert(TheCall->getNumArgs() == 1 && "expected 1 arg");
3829 // Update return type to be the attributed resource type from arg0.
3830 QualType ResourceTy = TheCall->getArg(0)->getType();
3831 TheCall->setType(ResourceTy);
3832 break;
3833 }
3834 case Builtin::BI__builtin_hlsl_resource_handlefrombinding: {
3835 assert(TheCall->getNumArgs() == 6 && "expected 6 args");
3836 // Update return type to be the attributed resource type from arg0.
3837 QualType ResourceTy = TheCall->getArg(0)->getType();
3838 TheCall->setType(ResourceTy);
3839 break;
3840 }
3841 case Builtin::BI__builtin_hlsl_resource_handlefromimplicitbinding: {
3842 assert(TheCall->getNumArgs() == 6 && "expected 6 args");
3843 // Update return type to be the attributed resource type from arg0.
3844 QualType ResourceTy = TheCall->getArg(0)->getType();
3845 TheCall->setType(ResourceTy);
3846 break;
3847 }
3848 case Builtin::BI__builtin_hlsl_resource_counterhandlefromimplicitbinding: {
3849 assert(TheCall->getNumArgs() == 3 && "expected 3 args");
3850 ASTContext &AST = SemaRef.getASTContext();
3851 QualType MainHandleTy = TheCall->getArg(0)->getType();
3852 auto *MainResType = MainHandleTy->getAs<HLSLAttributedResourceType>();
3853 auto MainAttrs = MainResType->getAttrs();
3854 assert(!MainAttrs.IsCounter && "cannot create a counter from a counter");
3855 MainAttrs.IsCounter = true;
3856 QualType CounterHandleTy = AST.getHLSLAttributedResourceType(
3857 MainResType->getWrappedType(), MainResType->getContainedType(),
3858 MainAttrs);
3859 // Update return type to be the attributed resource type from arg0
3860 // with added IsCounter flag.
3861 TheCall->setType(CounterHandleTy);
3862 break;
3863 }
3864 case Builtin::BI__builtin_hlsl_and:
3865 case Builtin::BI__builtin_hlsl_or: {
3866 if (SemaRef.checkArgCount(TheCall, 2))
3867 return true;
3868 if (CheckScalarOrVectorOrMatrix(&SemaRef, TheCall, getASTContext().BoolTy,
3869 0))
3870 return true;
3871 if (CheckAllArgsHaveSameType(&SemaRef, TheCall))
3872 return true;
3873
3874 ExprResult A = TheCall->getArg(0);
3875 QualType ArgTyA = A.get()->getType();
3876 // return type is the same as the input type
3877 TheCall->setType(ArgTyA);
3878 break;
3879 }
3880 case Builtin::BI__builtin_hlsl_all:
3881 case Builtin::BI__builtin_hlsl_any: {
3882 if (SemaRef.checkArgCount(TheCall, 1))
3883 return true;
3884 if (CheckAnyScalarOrVector(&SemaRef, TheCall, 0))
3885 return true;
3886 break;
3887 }
3888 case Builtin::BI__builtin_hlsl_asdouble: {
3889 if (SemaRef.checkArgCount(TheCall, 2))
3890 return true;
3892 &SemaRef, TheCall,
3893 /*only check for uint*/ SemaRef.Context.UnsignedIntTy,
3894 /* arg index */ 0))
3895 return true;
3897 &SemaRef, TheCall,
3898 /*only check for uint*/ SemaRef.Context.UnsignedIntTy,
3899 /* arg index */ 1))
3900 return true;
3901 if (CheckAllArgsHaveSameType(&SemaRef, TheCall))
3902 return true;
3903
3904 SetElementTypeAsReturnType(&SemaRef, TheCall, getASTContext().DoubleTy);
3905 break;
3906 }
3907 case Builtin::BI__builtin_hlsl_elementwise_clamp: {
3908 if (SemaRef.BuiltinElementwiseTernaryMath(
3909 TheCall, /*ArgTyRestr=*/
3911 return true;
3912 break;
3913 }
3914 case Builtin::BI__builtin_hlsl_dot: {
3915 // arg count is checked by BuiltinVectorToScalarMath
3916 if (SemaRef.BuiltinVectorToScalarMath(TheCall))
3917 return true;
3919 return true;
3920 break;
3921 }
3922 case Builtin::BI__builtin_hlsl_elementwise_firstbithigh:
3923 case Builtin::BI__builtin_hlsl_elementwise_firstbitlow: {
3924 if (SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
3925 return true;
3926
3927 const Expr *Arg = TheCall->getArg(0);
3928 QualType ArgTy = Arg->getType();
3929 QualType EltTy = ArgTy;
3930
3931 QualType ResTy = SemaRef.Context.UnsignedIntTy;
3932
3933 if (auto *VecTy = EltTy->getAs<VectorType>()) {
3934 EltTy = VecTy->getElementType();
3935 ResTy = SemaRef.Context.getExtVectorType(ResTy, VecTy->getNumElements());
3936 }
3937
3938 if (!EltTy->isIntegerType()) {
3939 Diag(Arg->getBeginLoc(), diag::err_builtin_invalid_arg_type)
3940 << 1 << /* scalar or vector of */ 5 << /* integer ty */ 1
3941 << /* no fp */ 0 << ArgTy;
3942 return true;
3943 }
3944
3945 TheCall->setType(ResTy);
3946 break;
3947 }
3948 case Builtin::BI__builtin_hlsl_select: {
3949 if (SemaRef.checkArgCount(TheCall, 3))
3950 return true;
3951 if (CheckScalarOrVector(&SemaRef, TheCall, getASTContext().BoolTy, 0))
3952 return true;
3953 QualType ArgTy = TheCall->getArg(0)->getType();
3954 if (ArgTy->isBooleanType() && CheckBoolSelect(&SemaRef, TheCall))
3955 return true;
3956 auto *VTy = ArgTy->getAs<VectorType>();
3957 if (VTy && VTy->getElementType()->isBooleanType() &&
3958 CheckVectorSelect(&SemaRef, TheCall))
3959 return true;
3960 break;
3961 }
3962 case Builtin::BI__builtin_hlsl_elementwise_saturate:
3963 case Builtin::BI__builtin_hlsl_elementwise_rcp: {
3964 if (SemaRef.checkArgCount(TheCall, 1))
3965 return true;
3966 if (!TheCall->getArg(0)
3967 ->getType()
3968 ->hasFloatingRepresentation()) // half or float or double
3969 return SemaRef.Diag(TheCall->getArg(0)->getBeginLoc(),
3970 diag::err_builtin_invalid_arg_type)
3971 << /* ordinal */ 1 << /* scalar or vector */ 5 << /* no int */ 0
3972 << /* fp */ 1 << TheCall->getArg(0)->getType();
3973 if (SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
3974 return true;
3975 break;
3976 }
3977 case Builtin::BI__builtin_hlsl_elementwise_degrees:
3978 case Builtin::BI__builtin_hlsl_elementwise_radians:
3979 case Builtin::BI__builtin_hlsl_elementwise_rsqrt:
3980 case Builtin::BI__builtin_hlsl_elementwise_frac:
3981 case Builtin::BI__builtin_hlsl_elementwise_ddx_coarse:
3982 case Builtin::BI__builtin_hlsl_elementwise_ddy_coarse:
3983 case Builtin::BI__builtin_hlsl_elementwise_ddx_fine:
3984 case Builtin::BI__builtin_hlsl_elementwise_ddy_fine: {
3985 if (SemaRef.checkArgCount(TheCall, 1))
3986 return true;
3987 if (CheckAllArgTypesAreCorrect(&SemaRef, TheCall,
3989 return true;
3990 if (SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
3991 return true;
3992 break;
3993 }
3994 case Builtin::BI__builtin_hlsl_elementwise_isinf:
3995 case Builtin::BI__builtin_hlsl_elementwise_isnan: {
3996 if (SemaRef.checkArgCount(TheCall, 1))
3997 return true;
3998 if (CheckAllArgTypesAreCorrect(&SemaRef, TheCall,
4000 return true;
4001 if (SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
4002 return true;
4004 break;
4005 }
4006 case Builtin::BI__builtin_hlsl_lerp: {
4007 if (SemaRef.checkArgCount(TheCall, 3))
4008 return true;
4009 if (CheckAllArgTypesAreCorrect(&SemaRef, TheCall,
4011 return true;
4012 if (CheckAllArgsHaveSameType(&SemaRef, TheCall))
4013 return true;
4014 if (SemaRef.BuiltinElementwiseTernaryMath(TheCall))
4015 return true;
4016 break;
4017 }
4018 case Builtin::BI__builtin_hlsl_mad: {
4019 if (SemaRef.BuiltinElementwiseTernaryMath(
4020 TheCall, /*ArgTyRestr=*/
4022 return true;
4023 break;
4024 }
4025 case Builtin::BI__builtin_hlsl_mul: {
4026 if (SemaRef.checkArgCount(TheCall, 2))
4027 return true;
4028
4029 Expr *Arg0 = TheCall->getArg(0);
4030 Expr *Arg1 = TheCall->getArg(1);
4031 QualType Ty0 = Arg0->getType();
4032 QualType Ty1 = Arg1->getType();
4033
4034 auto getElemType = [](QualType T) -> QualType {
4035 if (const auto *VTy = T->getAs<VectorType>())
4036 return VTy->getElementType();
4037 if (const auto *MTy = T->getAs<ConstantMatrixType>())
4038 return MTy->getElementType();
4039 return T;
4040 };
4041
4042 QualType EltTy0 = getElemType(Ty0);
4043
4044 bool IsVec0 = Ty0->isVectorType();
4045 bool IsMat0 = Ty0->isConstantMatrixType();
4046 bool IsVec1 = Ty1->isVectorType();
4047 bool IsMat1 = Ty1->isConstantMatrixType();
4048
4049 QualType RetTy;
4050
4051 if (IsVec0 && IsMat1) {
4052 auto *MatTy = Ty1->castAs<ConstantMatrixType>();
4053 RetTy = getASTContext().getExtVectorType(EltTy0, MatTy->getNumColumns());
4054 } else if (IsMat0 && IsVec1) {
4055 auto *MatTy = Ty0->castAs<ConstantMatrixType>();
4056 RetTy = getASTContext().getExtVectorType(EltTy0, MatTy->getNumRows());
4057 } else {
4058 assert(IsMat0 && IsMat1);
4059 auto *MatTy0 = Ty0->castAs<ConstantMatrixType>();
4060 auto *MatTy1 = Ty1->castAs<ConstantMatrixType>();
4062 EltTy0, MatTy0->getNumRows(), MatTy1->getNumColumns());
4063 }
4064
4065 TheCall->setType(RetTy);
4066 break;
4067 }
4068 case Builtin::BI__builtin_hlsl_normalize: {
4069 if (SemaRef.checkArgCount(TheCall, 1))
4070 return true;
4071 if (CheckAllArgTypesAreCorrect(&SemaRef, TheCall,
4073 return true;
4074 ExprResult A = TheCall->getArg(0);
4075 QualType ArgTyA = A.get()->getType();
4076 // return type is the same as the input type
4077 TheCall->setType(ArgTyA);
4078 break;
4079 }
4080 case Builtin::BI__builtin_hlsl_transpose: {
4081 if (SemaRef.checkArgCount(TheCall, 1))
4082 return true;
4083
4084 Expr *Arg = TheCall->getArg(0);
4085 QualType ArgTy = Arg->getType();
4086
4087 const auto *MatTy = ArgTy->getAs<ConstantMatrixType>();
4088 if (!MatTy) {
4089 SemaRef.Diag(Arg->getBeginLoc(), diag::err_builtin_invalid_arg_type)
4090 << 1 << /* matrix */ 3 << /* no int */ 0 << /* no fp */ 0 << ArgTy;
4091 return true;
4092 }
4093
4095 MatTy->getElementType(), MatTy->getNumColumns(), MatTy->getNumRows());
4096 TheCall->setType(RetTy);
4097 break;
4098 }
4099 case Builtin::BI__builtin_hlsl_elementwise_sign: {
4100 if (SemaRef.PrepareBuiltinElementwiseMathOneArgCall(TheCall))
4101 return true;
4102 if (CheckAllArgTypesAreCorrect(&SemaRef, TheCall,
4104 return true;
4106 break;
4107 }
4108 case Builtin::BI__builtin_hlsl_step: {
4109 if (SemaRef.checkArgCount(TheCall, 2))
4110 return true;
4111 if (CheckAllArgTypesAreCorrect(&SemaRef, TheCall,
4113 return true;
4114
4115 ExprResult A = TheCall->getArg(0);
4116 QualType ArgTyA = A.get()->getType();
4117 // return type is the same as the input type
4118 TheCall->setType(ArgTyA);
4119 break;
4120 }
4121 case Builtin::BI__builtin_hlsl_wave_active_all_equal: {
4122 if (SemaRef.checkArgCount(TheCall, 1))
4123 return true;
4124
4125 // Ensure input expr type is a scalar/vector
4126 if (CheckAnyScalarOrVector(&SemaRef, TheCall, 0))
4127 return true;
4128
4129 QualType InputTy = TheCall->getArg(0)->getType();
4130 ASTContext &Ctx = getASTContext();
4131
4132 QualType RetTy;
4133
4134 // If vector, construct bool vector of same size
4135 if (const auto *VecTy = InputTy->getAs<ExtVectorType>()) {
4136 unsigned NumElts = VecTy->getNumElements();
4137 RetTy = Ctx.getExtVectorType(Ctx.BoolTy, NumElts);
4138 } else {
4139 // Scalar case
4140 RetTy = Ctx.BoolTy;
4141 }
4142
4143 TheCall->setType(RetTy);
4144 break;
4145 }
4146 case Builtin::BI__builtin_hlsl_wave_active_max:
4147 case Builtin::BI__builtin_hlsl_wave_active_min:
4148 case Builtin::BI__builtin_hlsl_wave_active_sum:
4149 case Builtin::BI__builtin_hlsl_wave_active_product: {
4150 if (SemaRef.checkArgCount(TheCall, 1))
4151 return true;
4152
4153 // Ensure input expr type is a scalar/vector and the same as the return type
4154 if (CheckAnyScalarOrVector(&SemaRef, TheCall, 0))
4155 return true;
4156 if (CheckWaveActive(&SemaRef, TheCall))
4157 return true;
4158 ExprResult Expr = TheCall->getArg(0);
4159 QualType ArgTyExpr = Expr.get()->getType();
4160 TheCall->setType(ArgTyExpr);
4161 break;
4162 }
4163 case Builtin::BI__builtin_hlsl_wave_active_bit_xor:
4164 case Builtin::BI__builtin_hlsl_wave_active_bit_or: {
4165 if (SemaRef.checkArgCount(TheCall, 1))
4166 return true;
4167
4168 // Ensure input expr type is a scalar/vector
4169 if (CheckAnyScalarOrVector(&SemaRef, TheCall, 0))
4170 return true;
4171
4172 if (CheckWaveActive(&SemaRef, TheCall))
4173 return true;
4174
4175 // Ensure the expr type is interpretable as a uint or vector<uint>
4176 ExprResult Expr = TheCall->getArg(0);
4177 QualType ArgTyExpr = Expr.get()->getType();
4178 auto *VTy = ArgTyExpr->getAs<VectorType>();
4179 if (!(ArgTyExpr->isIntegerType() ||
4180 (VTy && VTy->getElementType()->isIntegerType()))) {
4181 SemaRef.Diag(TheCall->getArg(0)->getBeginLoc(),
4182 diag::err_builtin_invalid_arg_type)
4183 << ArgTyExpr << SemaRef.Context.UnsignedIntTy << 1 << 0 << 0;
4184 return true;
4185 }
4186
4187 // Ensure input expr type is the same as the return type
4188 TheCall->setType(ArgTyExpr);
4189 break;
4190 }
4191 // Note these are llvm builtins that we want to catch invalid intrinsic
4192 // generation. Normal handling of these builtins will occur elsewhere.
4193 case Builtin::BI__builtin_elementwise_bitreverse: {
4194 // does not include a check for number of arguments
4195 // because that is done previously
4196 if (CheckAllArgTypesAreCorrect(&SemaRef, TheCall,
4198 return true;
4199 break;
4200 }
4201 case Builtin::BI__builtin_hlsl_wave_prefix_count_bits: {
4202 if (SemaRef.checkArgCount(TheCall, 1))
4203 return true;
4204
4205 QualType ArgType = TheCall->getArg(0)->getType();
4206
4207 if (!(ArgType->isScalarType())) {
4208 SemaRef.Diag(TheCall->getArg(0)->getBeginLoc(),
4209 diag::err_typecheck_expect_any_scalar_or_vector)
4210 << ArgType << 0;
4211 return true;
4212 }
4213
4214 if (!(ArgType->isBooleanType())) {
4215 SemaRef.Diag(TheCall->getArg(0)->getBeginLoc(),
4216 diag::err_typecheck_expect_any_scalar_or_vector)
4217 << ArgType << 0;
4218 return true;
4219 }
4220
4221 break;
4222 }
4223 case Builtin::BI__builtin_hlsl_wave_read_lane_at: {
4224 if (SemaRef.checkArgCount(TheCall, 2))
4225 return true;
4226
4227 // Ensure index parameter type can be interpreted as a uint
4228 ExprResult Index = TheCall->getArg(1);
4229 QualType ArgTyIndex = Index.get()->getType();
4230 if (!ArgTyIndex->isIntegerType()) {
4231 SemaRef.Diag(TheCall->getArg(1)->getBeginLoc(),
4232 diag::err_typecheck_convert_incompatible)
4233 << ArgTyIndex << SemaRef.Context.UnsignedIntTy << 1 << 0 << 0;
4234 return true;
4235 }
4236
4237 // Ensure input expr type is a scalar/vector and the same as the return type
4238 if (CheckAnyScalarOrVector(&SemaRef, TheCall, 0))
4239 return true;
4240
4241 ExprResult Expr = TheCall->getArg(0);
4242 QualType ArgTyExpr = Expr.get()->getType();
4243 TheCall->setType(ArgTyExpr);
4244 break;
4245 }
4246 case Builtin::BI__builtin_hlsl_wave_get_lane_index: {
4247 if (SemaRef.checkArgCount(TheCall, 0))
4248 return true;
4249 break;
4250 }
4251 case Builtin::BI__builtin_hlsl_wave_prefix_sum:
4252 case Builtin::BI__builtin_hlsl_wave_prefix_product: {
4253 if (SemaRef.checkArgCount(TheCall, 1))
4254 return true;
4255
4256 // Ensure input expr type is a scalar/vector and the same as the return type
4257 if (CheckAnyScalarOrVector(&SemaRef, TheCall, 0))
4258 return true;
4259 if (CheckWavePrefix(&SemaRef, TheCall))
4260 return true;
4261 ExprResult Expr = TheCall->getArg(0);
4262 QualType ArgTyExpr = Expr.get()->getType();
4263 TheCall->setType(ArgTyExpr);
4264 break;
4265 }
4266 case Builtin::BI__builtin_hlsl_quad_read_across_x: {
4267 if (SemaRef.checkArgCount(TheCall, 1))
4268 return true;
4269
4270 if (CheckAnyScalarOrVector(&SemaRef, TheCall, 0))
4271 return true;
4272 if (CheckNotBoolScalarOrVector(&SemaRef, TheCall, 0))
4273 return true;
4274 ExprResult Expr = TheCall->getArg(0);
4275 QualType ArgTyExpr = Expr.get()->getType();
4276 TheCall->setType(ArgTyExpr);
4277 break;
4278 }
4279 case Builtin::BI__builtin_hlsl_elementwise_splitdouble: {
4280 if (SemaRef.checkArgCount(TheCall, 3))
4281 return true;
4282
4283 if (CheckScalarOrVector(&SemaRef, TheCall, SemaRef.Context.DoubleTy, 0) ||
4284 CheckScalarOrVector(&SemaRef, TheCall, SemaRef.Context.UnsignedIntTy,
4285 1) ||
4286 CheckScalarOrVector(&SemaRef, TheCall, SemaRef.Context.UnsignedIntTy,
4287 2))
4288 return true;
4289
4290 if (CheckModifiableLValue(&SemaRef, TheCall, 1) ||
4291 CheckModifiableLValue(&SemaRef, TheCall, 2))
4292 return true;
4293 break;
4294 }
4295 case Builtin::BI__builtin_hlsl_elementwise_clip: {
4296 if (SemaRef.checkArgCount(TheCall, 1))
4297 return true;
4298
4299 if (CheckScalarOrVector(&SemaRef, TheCall, SemaRef.Context.FloatTy, 0))
4300 return true;
4301 break;
4302 }
4303 case Builtin::BI__builtin_elementwise_acos:
4304 case Builtin::BI__builtin_elementwise_asin:
4305 case Builtin::BI__builtin_elementwise_atan:
4306 case Builtin::BI__builtin_elementwise_atan2:
4307 case Builtin::BI__builtin_elementwise_ceil:
4308 case Builtin::BI__builtin_elementwise_cos:
4309 case Builtin::BI__builtin_elementwise_cosh:
4310 case Builtin::BI__builtin_elementwise_exp:
4311 case Builtin::BI__builtin_elementwise_exp2:
4312 case Builtin::BI__builtin_elementwise_exp10:
4313 case Builtin::BI__builtin_elementwise_floor:
4314 case Builtin::BI__builtin_elementwise_fmod:
4315 case Builtin::BI__builtin_elementwise_log:
4316 case Builtin::BI__builtin_elementwise_log2:
4317 case Builtin::BI__builtin_elementwise_log10:
4318 case Builtin::BI__builtin_elementwise_pow:
4319 case Builtin::BI__builtin_elementwise_roundeven:
4320 case Builtin::BI__builtin_elementwise_sin:
4321 case Builtin::BI__builtin_elementwise_sinh:
4322 case Builtin::BI__builtin_elementwise_sqrt:
4323 case Builtin::BI__builtin_elementwise_tan:
4324 case Builtin::BI__builtin_elementwise_tanh:
4325 case Builtin::BI__builtin_elementwise_trunc: {
4326 if (CheckAllArgTypesAreCorrect(&SemaRef, TheCall,
4328 return true;
4329 break;
4330 }
4331 case Builtin::BI__builtin_hlsl_buffer_update_counter: {
4332 assert(TheCall->getNumArgs() == 2 && "expected 2 args");
4333 auto checkResTy = [](const HLSLAttributedResourceType *ResTy) -> bool {
4334 return !(ResTy->getAttrs().ResourceClass == ResourceClass::UAV &&
4335 ResTy->getAttrs().RawBuffer && ResTy->hasContainedType());
4336 };
4337 if (CheckResourceHandle(&SemaRef, TheCall, 0, checkResTy))
4338 return true;
4339 Expr *OffsetExpr = TheCall->getArg(1);
4340 std::optional<llvm::APSInt> Offset =
4341 OffsetExpr->getIntegerConstantExpr(SemaRef.getASTContext());
4342 if (!Offset.has_value() || std::abs(Offset->getExtValue()) != 1) {
4343 SemaRef.Diag(TheCall->getArg(1)->getBeginLoc(),
4344 diag::err_hlsl_expect_arg_const_int_one_or_neg_one)
4345 << 1;
4346 return true;
4347 }
4348 break;
4349 }
4350 case Builtin::BI__builtin_hlsl_elementwise_f16tof32: {
4351 if (SemaRef.checkArgCount(TheCall, 1))
4352 return true;
4353 if (CheckAllArgTypesAreCorrect(&SemaRef, TheCall,
4355 return true;
4356 // ensure arg integers are 32 bits
4357 if (CheckExpectedBitWidth(&SemaRef, TheCall, 0, 32))
4358 return true;
4359 // check it wasn't a bool type
4360 QualType ArgTy = TheCall->getArg(0)->getType();
4361 if (auto *VTy = ArgTy->getAs<VectorType>())
4362 ArgTy = VTy->getElementType();
4363 if (ArgTy->isBooleanType()) {
4364 SemaRef.Diag(TheCall->getArg(0)->getBeginLoc(),
4365 diag::err_builtin_invalid_arg_type)
4366 << 1 << /* scalar or vector of */ 5 << /* unsigned int */ 3
4367 << /* no fp */ 0 << TheCall->getArg(0)->getType();
4368 return true;
4369 }
4370
4371 SetElementTypeAsReturnType(&SemaRef, TheCall, getASTContext().FloatTy);
4372 break;
4373 }
4374 case Builtin::BI__builtin_hlsl_elementwise_f32tof16: {
4375 if (SemaRef.checkArgCount(TheCall, 1))
4376 return true;
4378 return true;
4380 getASTContext().UnsignedIntTy);
4381 break;
4382 }
4383 }
4384 return false;
4385}
4386
4390 WorkList.push_back(BaseTy);
4391 while (!WorkList.empty()) {
4392 QualType T = WorkList.pop_back_val();
4393 T = T.getCanonicalType().getUnqualifiedType();
4394 if (const auto *AT = dyn_cast<ConstantArrayType>(T)) {
4395 llvm::SmallVector<QualType, 16> ElementFields;
4396 // Generally I've avoided recursion in this algorithm, but arrays of
4397 // structs could be time-consuming to flatten and churn through on the
4398 // work list. Hopefully nesting arrays of structs containing arrays
4399 // of structs too many levels deep is unlikely.
4400 BuildFlattenedTypeList(AT->getElementType(), ElementFields);
4401 // Repeat the element's field list n times.
4402 for (uint64_t Ct = 0; Ct < AT->getZExtSize(); ++Ct)
4403 llvm::append_range(List, ElementFields);
4404 continue;
4405 }
4406 // Vectors can only have element types that are builtin types, so this can
4407 // add directly to the list instead of to the WorkList.
4408 if (const auto *VT = dyn_cast<VectorType>(T)) {
4409 List.insert(List.end(), VT->getNumElements(), VT->getElementType());
4410 continue;
4411 }
4412 if (const auto *MT = dyn_cast<ConstantMatrixType>(T)) {
4413 List.insert(List.end(), MT->getNumElementsFlattened(),
4414 MT->getElementType());
4415 continue;
4416 }
4417 if (const auto *RD = T->getAsCXXRecordDecl()) {
4418 if (RD->isStandardLayout())
4419 RD = RD->getStandardLayoutBaseWithFields();
4420
4421 // For types that we shouldn't decompose (unions and non-aggregates), just
4422 // add the type itself to the list.
4423 if (RD->isUnion() || !RD->isAggregate()) {
4424 List.push_back(T);
4425 continue;
4426 }
4427
4429 for (const auto *FD : RD->fields())
4430 if (!FD->isUnnamedBitField())
4431 FieldTypes.push_back(FD->getType());
4432 // Reverse the newly added sub-range.
4433 std::reverse(FieldTypes.begin(), FieldTypes.end());
4434 llvm::append_range(WorkList, FieldTypes);
4435
4436 // If this wasn't a standard layout type we may also have some base
4437 // classes to deal with.
4438 if (!RD->isStandardLayout()) {
4439 FieldTypes.clear();
4440 for (const auto &Base : RD->bases())
4441 FieldTypes.push_back(Base.getType());
4442 std::reverse(FieldTypes.begin(), FieldTypes.end());
4443 llvm::append_range(WorkList, FieldTypes);
4444 }
4445 continue;
4446 }
4447 List.push_back(T);
4448 }
4449}
4450
4452 // null and array types are not allowed.
4453 if (QT.isNull() || QT->isArrayType())
4454 return false;
4455
4456 // UDT types are not allowed
4457 if (QT->isRecordType())
4458 return false;
4459
4460 if (QT->isBooleanType() || QT->isEnumeralType())
4461 return false;
4462
4463 // the only other valid builtin types are scalars or vectors
4464 if (QT->isArithmeticType()) {
4465 if (SemaRef.Context.getTypeSize(QT) / 8 > 16)
4466 return false;
4467 return true;
4468 }
4469
4470 if (const VectorType *VT = QT->getAs<VectorType>()) {
4471 int ArraySize = VT->getNumElements();
4472
4473 if (ArraySize > 4)
4474 return false;
4475
4476 QualType ElTy = VT->getElementType();
4477 if (ElTy->isBooleanType())
4478 return false;
4479
4480 if (SemaRef.Context.getTypeSize(QT) / 8 > 16)
4481 return false;
4482 return true;
4483 }
4484
4485 return false;
4486}
4487
4489 if (T1.isNull() || T2.isNull())
4490 return false;
4491
4494
4495 // If both types are the same canonical type, they're obviously compatible.
4496 if (SemaRef.getASTContext().hasSameType(T1, T2))
4497 return true;
4498
4500 BuildFlattenedTypeList(T1, T1Types);
4502 BuildFlattenedTypeList(T2, T2Types);
4503
4504 // Check the flattened type list
4505 return llvm::equal(T1Types, T2Types,
4506 [this](QualType LHS, QualType RHS) -> bool {
4507 return SemaRef.IsLayoutCompatible(LHS, RHS);
4508 });
4509}
4510
4512 FunctionDecl *Old) {
4513 if (New->getNumParams() != Old->getNumParams())
4514 return true;
4515
4516 bool HadError = false;
4517
4518 for (unsigned i = 0, e = New->getNumParams(); i != e; ++i) {
4519 ParmVarDecl *NewParam = New->getParamDecl(i);
4520 ParmVarDecl *OldParam = Old->getParamDecl(i);
4521
4522 // HLSL parameter declarations for inout and out must match between
4523 // declarations. In HLSL inout and out are ambiguous at the call site,
4524 // but have different calling behavior, so you cannot overload a
4525 // method based on a difference between inout and out annotations.
4526 const auto *NDAttr = NewParam->getAttr<HLSLParamModifierAttr>();
4527 unsigned NSpellingIdx = (NDAttr ? NDAttr->getSpellingListIndex() : 0);
4528 const auto *ODAttr = OldParam->getAttr<HLSLParamModifierAttr>();
4529 unsigned OSpellingIdx = (ODAttr ? ODAttr->getSpellingListIndex() : 0);
4530
4531 if (NSpellingIdx != OSpellingIdx) {
4532 SemaRef.Diag(NewParam->getLocation(),
4533 diag::err_hlsl_param_qualifier_mismatch)
4534 << NDAttr << NewParam;
4535 SemaRef.Diag(OldParam->getLocation(), diag::note_previous_declaration_as)
4536 << ODAttr;
4537 HadError = true;
4538 }
4539 }
4540 return HadError;
4541}
4542
4543// Generally follows PerformScalarCast, with cases reordered for
4544// clarity of what types are supported
4546
4547 if (!SrcTy->isScalarType() || !DestTy->isScalarType())
4548 return false;
4549
4550 if (SemaRef.getASTContext().hasSameUnqualifiedType(SrcTy, DestTy))
4551 return true;
4552
4553 switch (SrcTy->getScalarTypeKind()) {
4554 case Type::STK_Bool: // casting from bool is like casting from an integer
4555 case Type::STK_Integral:
4556 switch (DestTy->getScalarTypeKind()) {
4557 case Type::STK_Bool:
4558 case Type::STK_Integral:
4559 case Type::STK_Floating:
4560 return true;
4561 case Type::STK_CPointer:
4565 llvm_unreachable("HLSL doesn't support pointers.");
4568 llvm_unreachable("HLSL doesn't support complex types.");
4570 llvm_unreachable("HLSL doesn't support fixed point types.");
4571 }
4572 llvm_unreachable("Should have returned before this");
4573
4574 case Type::STK_Floating:
4575 switch (DestTy->getScalarTypeKind()) {
4576 case Type::STK_Floating:
4577 case Type::STK_Bool:
4578 case Type::STK_Integral:
4579 return true;
4582 llvm_unreachable("HLSL doesn't support complex types.");
4584 llvm_unreachable("HLSL doesn't support fixed point types.");
4585 case Type::STK_CPointer:
4589 llvm_unreachable("HLSL doesn't support pointers.");
4590 }
4591 llvm_unreachable("Should have returned before this");
4592
4594 case Type::STK_CPointer:
4597 llvm_unreachable("HLSL doesn't support pointers.");
4598
4600 llvm_unreachable("HLSL doesn't support fixed point types.");
4601
4604 llvm_unreachable("HLSL doesn't support complex types.");
4605 }
4606
4607 llvm_unreachable("Unhandled scalar cast");
4608}
4609
4610// Can perform an HLSL Aggregate splat cast if the Dest is an aggregate and the
4611// Src is a scalar or a vector of length 1
4612// Or if Dest is a vector and Src is a vector of length 1
4614
4615 QualType SrcTy = Src->getType();
4616 // Not a valid HLSL Aggregate Splat cast if Dest is a scalar or if this is
4617 // going to be a vector splat from a scalar.
4618 if ((SrcTy->isScalarType() && DestTy->isVectorType()) ||
4619 DestTy->isScalarType())
4620 return false;
4621
4622 const VectorType *SrcVecTy = SrcTy->getAs<VectorType>();
4623
4624 // Src isn't a scalar or a vector of length 1
4625 if (!SrcTy->isScalarType() && !(SrcVecTy && SrcVecTy->getNumElements() == 1))
4626 return false;
4627
4628 if (SrcVecTy)
4629 SrcTy = SrcVecTy->getElementType();
4630
4632 BuildFlattenedTypeList(DestTy, DestTypes);
4633
4634 for (unsigned I = 0, Size = DestTypes.size(); I < Size; ++I) {
4635 if (DestTypes[I]->isUnionType())
4636 return false;
4637 if (!CanPerformScalarCast(SrcTy, DestTypes[I]))
4638 return false;
4639 }
4640 return true;
4641}
4642
4643// Can we perform an HLSL Elementwise cast?
4645
4646 // Don't handle casts where LHS and RHS are any combination of scalar/vector
4647 // There must be an aggregate somewhere
4648 QualType SrcTy = Src->getType();
4649 if (SrcTy->isScalarType()) // always a splat and this cast doesn't handle that
4650 return false;
4651
4652 if (SrcTy->isVectorType() &&
4653 (DestTy->isScalarType() || DestTy->isVectorType()))
4654 return false;
4655
4656 if (SrcTy->isConstantMatrixType() &&
4657 (DestTy->isScalarType() || DestTy->isConstantMatrixType()))
4658 return false;
4659
4661 BuildFlattenedTypeList(DestTy, DestTypes);
4663 BuildFlattenedTypeList(SrcTy, SrcTypes);
4664
4665 // Usually the size of SrcTypes must be greater than or equal to the size of
4666 // DestTypes.
4667 if (SrcTypes.size() < DestTypes.size())
4668 return false;
4669
4670 unsigned SrcSize = SrcTypes.size();
4671 unsigned DstSize = DestTypes.size();
4672 unsigned I;
4673 for (I = 0; I < DstSize && I < SrcSize; I++) {
4674 if (SrcTypes[I]->isUnionType() || DestTypes[I]->isUnionType())
4675 return false;
4676 if (!CanPerformScalarCast(SrcTypes[I], DestTypes[I])) {
4677 return false;
4678 }
4679 }
4680
4681 // check the rest of the source type for unions.
4682 for (; I < SrcSize; I++) {
4683 if (SrcTypes[I]->isUnionType())
4684 return false;
4685 }
4686 return true;
4687}
4688
4690 assert(Param->hasAttr<HLSLParamModifierAttr>() &&
4691 "We should not get here without a parameter modifier expression");
4692 const auto *Attr = Param->getAttr<HLSLParamModifierAttr>();
4693 if (Attr->getABI() == ParameterABI::Ordinary)
4694 return ExprResult(Arg);
4695
4696 bool IsInOut = Attr->getABI() == ParameterABI::HLSLInOut;
4697 if (!Arg->isLValue()) {
4698 SemaRef.Diag(Arg->getBeginLoc(), diag::error_hlsl_inout_lvalue)
4699 << Arg << (IsInOut ? 1 : 0);
4700 return ExprError();
4701 }
4702
4703 ASTContext &Ctx = SemaRef.getASTContext();
4704
4705 QualType Ty = Param->getType().getNonLValueExprType(Ctx);
4706
4707 // HLSL allows implicit conversions from scalars to vectors, but not the
4708 // inverse, so we need to disallow `inout` with scalar->vector or
4709 // scalar->matrix conversions.
4710 if (Arg->getType()->isScalarType() != Ty->isScalarType()) {
4711 SemaRef.Diag(Arg->getBeginLoc(), diag::error_hlsl_inout_scalar_extension)
4712 << Arg << (IsInOut ? 1 : 0);
4713 return ExprError();
4714 }
4715
4716 auto *ArgOpV = new (Ctx) OpaqueValueExpr(Param->getBeginLoc(), Arg->getType(),
4717 VK_LValue, OK_Ordinary, Arg);
4718
4719 // Parameters are initialized via copy initialization. This allows for
4720 // overload resolution of argument constructors.
4721 InitializedEntity Entity =
4723 ExprResult Res =
4724 SemaRef.PerformCopyInitialization(Entity, Param->getBeginLoc(), ArgOpV);
4725 if (Res.isInvalid())
4726 return ExprError();
4727 Expr *Base = Res.get();
4728 // After the cast, drop the reference type when creating the exprs.
4729 Ty = Ty.getNonLValueExprType(Ctx);
4730 auto *OpV = new (Ctx)
4731 OpaqueValueExpr(Param->getBeginLoc(), Ty, VK_LValue, OK_Ordinary, Base);
4732
4733 // Writebacks are performed with `=` binary operator, which allows for
4734 // overload resolution on writeback result expressions.
4735 Res = SemaRef.ActOnBinOp(SemaRef.getCurScope(), Param->getBeginLoc(),
4736 tok::equal, ArgOpV, OpV);
4737
4738 if (Res.isInvalid())
4739 return ExprError();
4740 Expr *Writeback = Res.get();
4741 auto *OutExpr =
4742 HLSLOutArgExpr::Create(Ctx, Ty, ArgOpV, OpV, Writeback, IsInOut);
4743
4744 return ExprResult(OutExpr);
4745}
4746
4748 // If HLSL gains support for references, all the cites that use this will need
4749 // to be updated with semantic checking to produce errors for
4750 // pointers/references.
4751 assert(!Ty->isReferenceType() &&
4752 "Pointer and reference types cannot be inout or out parameters");
4753 Ty = SemaRef.getASTContext().getLValueReferenceType(Ty);
4754 Ty.addRestrict();
4755 return Ty;
4756}
4757
4758// Returns true if the type has a non-empty constant buffer layout (if it is
4759// scalar, vector or matrix, or if it contains any of these.
4761 const Type *Ty = QT->getUnqualifiedDesugaredType();
4762 if (Ty->isScalarType() || Ty->isVectorType() || Ty->isMatrixType())
4763 return true;
4764
4766 return false;
4767
4768 if (const auto *RD = Ty->getAsCXXRecordDecl()) {
4769 for (const auto *FD : RD->fields()) {
4771 return true;
4772 }
4773 assert(RD->getNumBases() <= 1 &&
4774 "HLSL doesn't support multiple inheritance");
4775 return RD->getNumBases()
4776 ? hasConstantBufferLayout(RD->bases_begin()->getType())
4777 : false;
4778 }
4779
4780 if (const auto *AT = dyn_cast<ArrayType>(Ty)) {
4781 if (const auto *CAT = dyn_cast<ConstantArrayType>(AT))
4782 if (isZeroSizedArray(CAT))
4783 return false;
4785 }
4786
4787 return false;
4788}
4789
4790static bool IsDefaultBufferConstantDecl(const ASTContext &Ctx, VarDecl *VD) {
4791 bool IsVulkan =
4792 Ctx.getTargetInfo().getTriple().getOS() == llvm::Triple::Vulkan;
4793 bool IsVKPushConstant = IsVulkan && VD->hasAttr<HLSLVkPushConstantAttr>();
4794 QualType QT = VD->getType();
4795 return VD->getDeclContext()->isTranslationUnit() &&
4796 QT.getAddressSpace() == LangAS::Default &&
4797 VD->getStorageClass() != SC_Static &&
4798 !VD->hasAttr<HLSLVkConstantIdAttr>() && !IsVKPushConstant &&
4800}
4801
4803 // The variable already has an address space (groupshared for ex).
4804 if (Decl->getType().hasAddressSpace())
4805 return;
4806
4807 if (Decl->getType()->isDependentType())
4808 return;
4809
4810 QualType Type = Decl->getType();
4811
4812 if (Decl->hasAttr<HLSLVkExtBuiltinInputAttr>()) {
4813 LangAS ImplAS = LangAS::hlsl_input;
4814 Type = SemaRef.getASTContext().getAddrSpaceQualType(Type, ImplAS);
4815 Decl->setType(Type);
4816 return;
4817 }
4818
4819 bool IsVulkan = getASTContext().getTargetInfo().getTriple().getOS() ==
4820 llvm::Triple::Vulkan;
4821 if (IsVulkan && Decl->hasAttr<HLSLVkPushConstantAttr>()) {
4822 if (HasDeclaredAPushConstant)
4823 SemaRef.Diag(Decl->getLocation(), diag::err_hlsl_push_constant_unique);
4824
4826 Type = SemaRef.getASTContext().getAddrSpaceQualType(Type, ImplAS);
4827 Decl->setType(Type);
4828 HasDeclaredAPushConstant = true;
4829 return;
4830 }
4831
4832 if (Type->isSamplerT() || Type->isVoidType())
4833 return;
4834
4835 // Resource handles.
4837 return;
4838
4839 // Only static globals belong to the Private address space.
4840 // Non-static globals belongs to the cbuffer.
4841 if (Decl->getStorageClass() != SC_Static && !Decl->isStaticDataMember())
4842 return;
4843
4845 Type = SemaRef.getASTContext().getAddrSpaceQualType(Type, ImplAS);
4846 Decl->setType(Type);
4847}
4848
4849namespace {
4850
4851// Helper class for assigning bindings to resources declared within a struct.
4852// It keeps track of all binding attributes declared on a struct instance, and
4853// the offsets for each register type that have been assigned so far.
4854// Handles both explicit and implicit bindings.
4855class StructBindingContext {
4856 // Bindings and offsets per register type. We only need to support four
4857 // register types - SRV (u), UAV (t), CBuffer (c), and Sampler (s).
4858 HLSLResourceBindingAttr *RegBindingsAttrs[4];
4859 unsigned RegBindingOffset[4];
4860
4861 // Make sure the RegisterType values are what we expect
4862 static_assert(static_cast<unsigned>(RegisterType::SRV) == 0 &&
4863 static_cast<unsigned>(RegisterType::UAV) == 1 &&
4864 static_cast<unsigned>(RegisterType::CBuffer) == 2 &&
4865 static_cast<unsigned>(RegisterType::Sampler) == 3,
4866 "unexpected register type values");
4867
4868 // Vulkan binding attribute does not vary by register type.
4869 HLSLVkBindingAttr *VkBindingAttr;
4870 unsigned VkBindingOffset;
4871
4872public:
4873 // Constructor: gather all binding attributes on a struct instance and
4874 // initialize offsets.
4875 StructBindingContext(VarDecl *VD) {
4876 for (unsigned i = 0; i < 4; ++i) {
4877 RegBindingsAttrs[i] = nullptr;
4878 RegBindingOffset[i] = 0;
4879 }
4880 VkBindingAttr = nullptr;
4881 VkBindingOffset = 0;
4882
4883 ASTContext &AST = VD->getASTContext();
4884 bool IsSpirv = AST.getTargetInfo().getTriple().isSPIRV();
4885
4886 for (Attr *A : VD->attrs()) {
4887 if (auto *RBA = dyn_cast<HLSLResourceBindingAttr>(A)) {
4888 RegisterType RegType = RBA->getRegisterType();
4889 unsigned RegTypeIdx = static_cast<unsigned>(RegType);
4890 // Ignore unsupported register annotations, such as 'c' or 'i'.
4891 if (RegTypeIdx < 4)
4892 RegBindingsAttrs[RegTypeIdx] = RBA;
4893 continue;
4894 }
4895 // Gather the Vulkan binding attributes only if the target is SPIR-V.
4896 if (IsSpirv) {
4897 if (auto *VBA = dyn_cast<HLSLVkBindingAttr>(A))
4898 VkBindingAttr = VBA;
4899 }
4900 }
4901 }
4902
4903 // Creates a binding attribute for a resource based on the gathered attributes
4904 // and the required register type and range.
4905 Attr *createBindingAttr(SemaHLSL &S, ASTContext &AST, RegisterType RegType,
4906 unsigned Range) {
4907 assert(static_cast<unsigned>(RegType) < 4 && "unexpected register type");
4908
4909 if (VkBindingAttr) {
4910 unsigned Offset = VkBindingOffset;
4911 VkBindingOffset += Range;
4912 return HLSLVkBindingAttr::CreateImplicit(
4913 AST, VkBindingAttr->getBinding() + Offset, VkBindingAttr->getSet(),
4914 VkBindingAttr->getRange());
4915 }
4916
4917 HLSLResourceBindingAttr *RBA =
4918 RegBindingsAttrs[static_cast<unsigned>(RegType)];
4919 HLSLResourceBindingAttr *NewAttr = nullptr;
4920
4921 if (RBA && RBA->hasRegisterSlot()) {
4922 // Explicit binding - create a new attribute with offseted slot number
4923 // based on the required register type.
4924 unsigned Offset = RegBindingOffset[static_cast<unsigned>(RegType)];
4925 RegBindingOffset[static_cast<unsigned>(RegType)] += Range;
4926
4927 unsigned NewSlotNumber = RBA->getSlotNumber() + Offset;
4928 StringRef NewSlotNumberStr =
4929 createRegisterString(AST, RBA->getRegisterType(), NewSlotNumber);
4930 NewAttr = HLSLResourceBindingAttr::CreateImplicit(
4931 AST, NewSlotNumberStr, RBA->getSpace(), RBA->getRange());
4932 NewAttr->setBinding(RegType, NewSlotNumber, RBA->getSpaceNumber());
4933 } else {
4934 // No binding attribute or space-only binding - create a binding
4935 // attribute for implicit binding.
4936 NewAttr = HLSLResourceBindingAttr::CreateImplicit(AST, "", "0", {});
4937 NewAttr->setBinding(RegType, std::nullopt,
4938 RBA ? RBA->getSpaceNumber() : 0);
4939 NewAttr->setImplicitBindingOrderID(S.getNextImplicitBindingOrderID());
4940 }
4941 return NewAttr;
4942 }
4943};
4944
4945// Creates a global variable declaration for a resource field embedded in a
4946// struct, assigns it a binding, initializes it, and associates it with the
4947// struct declaration via an HLSLAssociatedResourceDeclAttr.
4948static void createGlobalResourceDeclForStruct(
4949 Sema &S, VarDecl *ParentVD, SourceLocation Loc, IdentifierInfo *Id,
4950 QualType ResTy, StructBindingContext &BindingCtx) {
4951 assert(isResourceRecordTypeOrArrayOf(ResTy) &&
4952 "expected resource type or array of resources");
4953
4954 DeclContext *DC = ParentVD->getNonTransparentDeclContext();
4955 assert(DC->isTranslationUnit() && "expected translation unit decl context");
4956
4957 ASTContext &AST = S.getASTContext();
4958 VarDecl *ResDecl =
4959 VarDecl::Create(AST, DC, Loc, Loc, Id, ResTy, nullptr, SC_None);
4960
4961 unsigned Range = 1;
4962 const HLSLAttributedResourceType *ResHandleTy = nullptr;
4963 if (const auto *AT = dyn_cast<ArrayType>(ResTy.getTypePtr())) {
4964 const auto *CAT = dyn_cast<ConstantArrayType>(AT);
4965 Range = CAT ? CAT->getSize().getZExtValue() : 0;
4966 ResHandleTy = getResourceArrayHandleType(ResTy);
4967 } else {
4968 ResHandleTy = HLSLAttributedResourceType::findHandleTypeOnResource(
4969 ResTy.getTypePtr());
4970 }
4971 // Add a binding attribute to the global resource declaration.
4972 Attr *BindingAttr = BindingCtx.createBindingAttr(
4973 S.HLSL(), AST, getRegisterType(ResHandleTy), Range);
4974 ResDecl->addAttr(BindingAttr);
4975 ResDecl->addAttr(InternalLinkageAttr::CreateImplicit(AST));
4976 ResDecl->setImplicit();
4977
4978 if (Range == 1)
4979 S.HLSL().initGlobalResourceDecl(ResDecl);
4980 else
4981 S.HLSL().initGlobalResourceArrayDecl(ResDecl);
4982
4983 ParentVD->addAttr(
4984 HLSLAssociatedResourceDeclAttr::CreateImplicit(AST, ResDecl));
4985 DC->addDecl(ResDecl);
4986
4987 DeclGroupRef DG(ResDecl);
4989}
4990
4991static void handleArrayOfStructWithResources(
4992 Sema &S, VarDecl *ParentVD, const ConstantArrayType *CAT,
4993 EmbeddedResourceNameBuilder &NameBuilder, StructBindingContext &BindingCtx);
4994
4995// Scans base and all fields of a struct/class type to find all embedded
4996// resources or resource arrays. Creates a global variable for each resource
4997// found.
4998static void handleStructWithResources(Sema &S, VarDecl *ParentVD,
4999 const CXXRecordDecl *RD,
5000 EmbeddedResourceNameBuilder &NameBuilder,
5001 StructBindingContext &BindingCtx) {
5002
5003 // Scan the base classes.
5004 assert(RD->getNumBases() <= 1 && "HLSL doesn't support multiple inheritance");
5005 const auto *BasesIt = RD->bases_begin();
5006 if (BasesIt != RD->bases_end()) {
5007 QualType QT = BasesIt->getType();
5008 if (QT->isHLSLIntangibleType()) {
5009 CXXRecordDecl *BaseRD = QT->getAsCXXRecordDecl();
5010 NameBuilder.pushBaseName(BaseRD->getName());
5011 handleStructWithResources(S, ParentVD, BaseRD, NameBuilder, BindingCtx);
5012 NameBuilder.pop();
5013 }
5014 }
5015 // Process this class fields.
5016 for (const FieldDecl *FD : RD->fields()) {
5017 QualType FDTy = FD->getType().getCanonicalType();
5018 if (!FDTy->isHLSLIntangibleType())
5019 continue;
5020
5021 NameBuilder.pushName(FD->getName());
5022
5024 IdentifierInfo *II = NameBuilder.getNameAsIdentifier(S.getASTContext());
5025 createGlobalResourceDeclForStruct(S, ParentVD, FD->getLocation(), II,
5026 FDTy, BindingCtx);
5027 } else if (const auto *RD = FDTy->getAsCXXRecordDecl()) {
5028 handleStructWithResources(S, ParentVD, RD, NameBuilder, BindingCtx);
5029
5030 } else if (const auto *ArrayTy = dyn_cast<ConstantArrayType>(FDTy)) {
5031 assert(!FDTy->isHLSLResourceRecordArray() &&
5032 "resource arrays should have been already handled");
5033 handleArrayOfStructWithResources(S, ParentVD, ArrayTy, NameBuilder,
5034 BindingCtx);
5035 }
5036 NameBuilder.pop();
5037 }
5038}
5039
5040// Processes array of structs with resources.
5041static void
5042handleArrayOfStructWithResources(Sema &S, VarDecl *ParentVD,
5043 const ConstantArrayType *CAT,
5044 EmbeddedResourceNameBuilder &NameBuilder,
5045 StructBindingContext &BindingCtx) {
5046
5047 QualType ElementTy = CAT->getElementType().getCanonicalType();
5048 assert(ElementTy->isHLSLIntangibleType() && "Expected HLSL intangible type");
5049
5050 const ConstantArrayType *SubCAT = dyn_cast<ConstantArrayType>(ElementTy);
5051 const CXXRecordDecl *ElementRD = ElementTy->getAsCXXRecordDecl();
5052
5053 if (!SubCAT && !ElementRD)
5054 return;
5055
5056 for (unsigned I = 0, E = CAT->getSize().getZExtValue(); I < E; ++I) {
5057 NameBuilder.pushArrayIndex(I);
5058 if (ElementRD)
5059 handleStructWithResources(S, ParentVD, ElementRD, NameBuilder,
5060 BindingCtx);
5061 else
5062 handleArrayOfStructWithResources(S, ParentVD, SubCAT, NameBuilder,
5063 BindingCtx);
5064 NameBuilder.pop();
5065 }
5066}
5067
5068} // namespace
5069
5070// Scans all fields of a user-defined struct (or array of structs)
5071// to find all embedded resources or resource arrays. For each resource
5072// a global variable of the resource type is created and associated
5073// with the parent declaration (VD) through a HLSLAssociatedResourceDeclAttr
5074// attribute.
5075void SemaHLSL::handleGlobalStructOrArrayOfWithResources(VarDecl *VD) {
5076 EmbeddedResourceNameBuilder NameBuilder(VD->getName());
5077 StructBindingContext BindingCtx(VD);
5078
5079 const Type *VDTy = VD->getType().getTypePtr();
5080 assert(VDTy->isHLSLIntangibleType() && !isResourceRecordTypeOrArrayOf(VD) &&
5081 "Expected non-resource struct or array type");
5082
5083 if (const CXXRecordDecl *RD = VDTy->getAsCXXRecordDecl()) {
5084 handleStructWithResources(SemaRef, VD, RD, NameBuilder, BindingCtx);
5085 return;
5086 }
5087
5088 if (const auto *CAT = dyn_cast<ConstantArrayType>(VDTy)) {
5089 handleArrayOfStructWithResources(SemaRef, VD, CAT, NameBuilder, BindingCtx);
5090 return;
5091 }
5092}
5093
5095 if (VD->hasGlobalStorage()) {
5096 // make sure the declaration has a complete type
5097 if (SemaRef.RequireCompleteType(
5098 VD->getLocation(),
5099 SemaRef.getASTContext().getBaseElementType(VD->getType()),
5100 diag::err_typecheck_decl_incomplete_type)) {
5101 VD->setInvalidDecl();
5103 return;
5104 }
5105
5106 // Global variables outside a cbuffer block that are not a resource, static,
5107 // groupshared, or an empty array or struct belong to the default constant
5108 // buffer $Globals (to be created at the end of the translation unit).
5110 // update address space to hlsl_constant
5113 VD->setType(NewTy);
5114 DefaultCBufferDecls.push_back(VD);
5115 }
5116
5117 // find all resources bindings on decl
5118 if (VD->getType()->isHLSLIntangibleType())
5119 collectResourceBindingsOnVarDecl(VD);
5120
5121 if (VD->hasAttr<HLSLVkConstantIdAttr>())
5123
5125 VD->getStorageClass() != SC_Static) {
5126 // Add internal linkage attribute to non-static resource variables. The
5127 // global externally visible storage is accessed through the handle, which
5128 // is a member. The variable itself is not externally visible.
5129 VD->addAttr(InternalLinkageAttr::CreateImplicit(getASTContext()));
5130 }
5131
5132 // process explicit bindings
5133 processExplicitBindingsOnDecl(VD);
5134
5135 // Add implicit binding attribute to non-static resource arrays.
5136 if (VD->getType()->isHLSLResourceRecordArray() &&
5137 VD->getStorageClass() != SC_Static) {
5138 // If the resource array does not have an explicit binding attribute,
5139 // create an implicit one. It will be used to transfer implicit binding
5140 // order_ID to codegen.
5141 ResourceBindingAttrs Binding(VD);
5142 if (!Binding.isExplicit()) {
5143 uint32_t OrderID = getNextImplicitBindingOrderID();
5144 if (Binding.hasBinding())
5145 Binding.setImplicitOrderID(OrderID);
5146 else {
5149 OrderID);
5150 // Re-create the binding object to pick up the new attribute.
5151 Binding = ResourceBindingAttrs(VD);
5152 }
5153 }
5154
5155 // Get to the base type of a potentially multi-dimensional array.
5157
5158 const CXXRecordDecl *RD = Ty->getAsCXXRecordDecl();
5159 if (hasCounterHandle(RD)) {
5160 if (!Binding.hasCounterImplicitOrderID()) {
5161 uint32_t OrderID = getNextImplicitBindingOrderID();
5162 Binding.setCounterImplicitOrderID(OrderID);
5163 }
5164 }
5165 }
5166
5167 // Process resources in user-defined structs, or arrays of such structs.
5168 const Type *VDTy = VD->getType().getTypePtr();
5169 if (VD->getStorageClass() != SC_Static && VDTy->isHLSLIntangibleType() &&
5171 handleGlobalStructOrArrayOfWithResources(VD);
5172
5173 // Mark groupshared variables as extern so they will have
5174 // external storage and won't be default initialized
5175 if (VD->hasAttr<HLSLGroupSharedAddressSpaceAttr>())
5177 }
5178
5180}
5181
5183 assert(VD->getType()->isHLSLResourceRecord() &&
5184 "expected resource record type");
5185
5186 ASTContext &AST = SemaRef.getASTContext();
5187 uint64_t UIntTySize = AST.getTypeSize(AST.UnsignedIntTy);
5188 uint64_t IntTySize = AST.getTypeSize(AST.IntTy);
5189
5190 // Gather resource binding attributes.
5191 ResourceBindingAttrs Binding(VD);
5192
5193 // Find correct initialization method and create its arguments.
5194 QualType ResourceTy = VD->getType();
5195 CXXRecordDecl *ResourceDecl = ResourceTy->getAsCXXRecordDecl();
5196 CXXMethodDecl *CreateMethod = nullptr;
5198
5199 bool HasCounter = hasCounterHandle(ResourceDecl);
5200 const char *CreateMethodName;
5201 if (Binding.isExplicit())
5202 CreateMethodName = HasCounter ? "__createFromBindingWithImplicitCounter"
5203 : "__createFromBinding";
5204 else
5205 CreateMethodName = HasCounter
5206 ? "__createFromImplicitBindingWithImplicitCounter"
5207 : "__createFromImplicitBinding";
5208
5209 CreateMethod =
5210 lookupMethod(SemaRef, ResourceDecl, CreateMethodName, VD->getLocation());
5211
5212 if (!CreateMethod)
5213 // This can happen if someone creates a struct that looks like an HLSL
5214 // resource record but does not have the required static create method.
5215 // No binding will be generated for it.
5216 return false;
5217
5218 if (Binding.isExplicit()) {
5219 IntegerLiteral *RegSlot =
5220 IntegerLiteral::Create(AST, llvm::APInt(UIntTySize, Binding.getSlot()),
5222 Args.push_back(RegSlot);
5223 } else {
5224 uint32_t OrderID = (Binding.hasImplicitOrderID())
5225 ? Binding.getImplicitOrderID()
5227 IntegerLiteral *OrderId =
5228 IntegerLiteral::Create(AST, llvm::APInt(UIntTySize, OrderID),
5230 Args.push_back(OrderId);
5231 }
5232
5233 IntegerLiteral *Space =
5234 IntegerLiteral::Create(AST, llvm::APInt(UIntTySize, Binding.getSpace()),
5236 Args.push_back(Space);
5237
5239 AST, llvm::APInt(IntTySize, 1), AST.IntTy, SourceLocation());
5240 Args.push_back(RangeSize);
5241
5243 AST, llvm::APInt(UIntTySize, 0), AST.UnsignedIntTy, SourceLocation());
5244 Args.push_back(Index);
5245
5246 StringRef VarName = VD->getName();
5248 AST, VarName, StringLiteralKind::Ordinary, false,
5249 AST.getStringLiteralArrayType(AST.CharTy.withConst(), VarName.size()),
5250 SourceLocation());
5252 AST, AST.getPointerType(AST.CharTy.withConst()), CK_ArrayToPointerDecay,
5253 Name, nullptr, VK_PRValue, FPOptionsOverride());
5254 Args.push_back(NameCast);
5255
5256 if (HasCounter) {
5257 // Will this be in the correct order?
5258 uint32_t CounterOrderID = getNextImplicitBindingOrderID();
5259 IntegerLiteral *CounterId =
5260 IntegerLiteral::Create(AST, llvm::APInt(UIntTySize, CounterOrderID),
5262 Args.push_back(CounterId);
5263 }
5264
5265 // Make sure the create method template is instantiated and emitted.
5266 if (!CreateMethod->isDefined() && CreateMethod->isTemplateInstantiation())
5267 SemaRef.InstantiateFunctionDefinition(VD->getLocation(), CreateMethod,
5268 true);
5269
5270 // Create CallExpr with a call to the static method and set it as the decl
5271 // initialization.
5273 AST, NestedNameSpecifierLoc(), SourceLocation(), CreateMethod, false,
5274 CreateMethod->getNameInfo(), CreateMethod->getType(), VK_PRValue);
5275
5276 auto *ImpCast = ImplicitCastExpr::Create(
5277 AST, AST.getPointerType(CreateMethod->getType()),
5278 CK_FunctionToPointerDecay, DRE, nullptr, VK_PRValue, FPOptionsOverride());
5279
5280 CallExpr *InitExpr =
5281 CallExpr::Create(AST, ImpCast, Args, ResourceTy, VK_PRValue,
5283 VD->setInit(InitExpr);
5285 SemaRef.CheckCompleteVariableDeclaration(VD);
5286 return true;
5287}
5288
5290 assert(VD->getType()->isHLSLResourceRecordArray() &&
5291 "expected array of resource records");
5292
5293 // Individual resources in a resource array are not initialized here. They
5294 // are initialized later on during codegen when the individual resources are
5295 // accessed. Codegen will emit a call to the resource initialization method
5296 // with the specified array index. We need to make sure though that the method
5297 // for the specific resource type is instantiated, so codegen can emit a call
5298 // to it when the array element is accessed.
5299
5300 // Find correct initialization method based on the resource binding
5301 // information.
5302 ASTContext &AST = SemaRef.getASTContext();
5303 QualType ResElementTy = AST.getBaseElementType(VD->getType());
5304 CXXRecordDecl *ResourceDecl = ResElementTy->getAsCXXRecordDecl();
5305 CXXMethodDecl *CreateMethod = nullptr;
5306
5307 bool HasCounter = hasCounterHandle(ResourceDecl);
5308 ResourceBindingAttrs ResourceAttrs(VD);
5309 if (ResourceAttrs.isExplicit())
5310 // Resource has explicit binding.
5311 CreateMethod =
5312 lookupMethod(SemaRef, ResourceDecl,
5313 HasCounter ? "__createFromBindingWithImplicitCounter"
5314 : "__createFromBinding",
5315 VD->getLocation());
5316 else
5317 // Resource has implicit binding.
5318 CreateMethod = lookupMethod(
5319 SemaRef, ResourceDecl,
5320 HasCounter ? "__createFromImplicitBindingWithImplicitCounter"
5321 : "__createFromImplicitBinding",
5322 VD->getLocation());
5323
5324 if (!CreateMethod)
5325 return false;
5326
5327 // Make sure the create method template is instantiated and emitted.
5328 if (!CreateMethod->isDefined() && CreateMethod->isTemplateInstantiation())
5329 SemaRef.InstantiateFunctionDefinition(VD->getLocation(), CreateMethod,
5330 true);
5331 return true;
5332}
5333
5334// Returns true if the initialization has been handled.
5335// Returns false to use default initialization.
5337 // Objects in the hlsl_constant address space are initialized
5338 // externally, so don't synthesize an implicit initializer.
5340 return true;
5341
5342 // Initialize non-static resources at the global scope.
5343 if (VD->hasGlobalStorage() && VD->getStorageClass() != SC_Static) {
5344 const Type *Ty = VD->getType().getTypePtr();
5345 if (Ty->isHLSLResourceRecord())
5346 return initGlobalResourceDecl(VD);
5347 if (Ty->isHLSLResourceRecordArray())
5348 return initGlobalResourceArrayDecl(VD);
5349 }
5350 return false;
5351}
5352
5353std::optional<const DeclBindingInfo *> SemaHLSL::inferGlobalBinding(Expr *E) {
5354 if (auto *Ternary = dyn_cast<ConditionalOperator>(E)) {
5355 auto TrueInfo = inferGlobalBinding(Ternary->getTrueExpr());
5356 auto FalseInfo = inferGlobalBinding(Ternary->getFalseExpr());
5357 if (!TrueInfo || !FalseInfo)
5358 return std::nullopt;
5359 if (*TrueInfo != *FalseInfo)
5360 return std::nullopt;
5361 return TrueInfo;
5362 }
5363
5364 if (auto *ASE = dyn_cast<ArraySubscriptExpr>(E))
5365 E = ASE->getBase()->IgnoreParenImpCasts();
5366
5367 if (DeclRefExpr *DRE = dyn_cast<DeclRefExpr>(E->IgnoreParens()))
5368 if (VarDecl *VD = dyn_cast<VarDecl>(DRE->getDecl())) {
5369 const Type *Ty = VD->getType()->getUnqualifiedDesugaredType();
5370 if (Ty->isArrayType())
5372
5373 if (const auto *AttrResType =
5374 HLSLAttributedResourceType::findHandleTypeOnResource(Ty)) {
5375 ResourceClass RC = AttrResType->getAttrs().ResourceClass;
5376 return Bindings.getDeclBindingInfo(VD, RC);
5377 }
5378 }
5379
5380 return nullptr;
5381}
5382
5383void SemaHLSL::trackLocalResource(VarDecl *VD, Expr *E) {
5384 std::optional<const DeclBindingInfo *> ExprBinding = inferGlobalBinding(E);
5385 if (!ExprBinding) {
5386 SemaRef.Diag(E->getBeginLoc(),
5387 diag::warn_hlsl_assigning_local_resource_is_not_unique)
5388 << E << VD;
5389 return; // Expr use multiple resources
5390 }
5391
5392 if (*ExprBinding == nullptr)
5393 return; // No binding could be inferred to track, return without error
5394
5395 auto PrevBinding = Assigns.find(VD);
5396 if (PrevBinding == Assigns.end()) {
5397 // No previous binding recorded, simply record the new assignment
5398 Assigns.insert({VD, *ExprBinding});
5399 return;
5400 }
5401
5402 // Otherwise, warn if the assignment implies different resource bindings
5403 if (*ExprBinding != PrevBinding->second) {
5404 SemaRef.Diag(E->getBeginLoc(),
5405 diag::warn_hlsl_assigning_local_resource_is_not_unique)
5406 << E << VD;
5407 SemaRef.Diag(VD->getLocation(), diag::note_var_declared_here) << VD;
5408 return;
5409 }
5410
5411 return;
5412}
5413
5415 Expr *RHSExpr, SourceLocation Loc) {
5416 assert((LHSExpr->getType()->isHLSLResourceRecord() ||
5417 LHSExpr->getType()->isHLSLResourceRecordArray()) &&
5418 "expected LHS to be a resource record or array of resource records");
5419 if (Opc != BO_Assign)
5420 return true;
5421
5422 // If LHS is an array subscript, get the underlying declaration.
5423 Expr *E = LHSExpr;
5424 while (auto *ASE = dyn_cast<ArraySubscriptExpr>(E))
5425 E = ASE->getBase()->IgnoreParenImpCasts();
5426
5427 // Report error if LHS is a non-static resource declared at a global scope.
5428 if (DeclRefExpr *DRE = dyn_cast<DeclRefExpr>(E->IgnoreParens())) {
5429 if (VarDecl *VD = dyn_cast<VarDecl>(DRE->getDecl())) {
5430 if (VD->hasGlobalStorage() && VD->getStorageClass() != SC_Static) {
5431 // assignment to global resource is not allowed
5432 SemaRef.Diag(Loc, diag::err_hlsl_assign_to_global_resource) << VD;
5433 SemaRef.Diag(VD->getLocation(), diag::note_var_declared_here) << VD;
5434 return false;
5435 }
5436
5437 trackLocalResource(VD, RHSExpr);
5438 }
5439 }
5440 return true;
5441}
5442
5443// Walks though the global variable declaration, collects all resource binding
5444// requirements and adds them to Bindings
5445void SemaHLSL::collectResourceBindingsOnVarDecl(VarDecl *VD) {
5446 assert(VD->hasGlobalStorage() && VD->getType()->isHLSLIntangibleType() &&
5447 "expected global variable that contains HLSL resource");
5448
5449 // Cbuffers and Tbuffers are HLSLBufferDecl types
5450 if (const HLSLBufferDecl *CBufferOrTBuffer = dyn_cast<HLSLBufferDecl>(VD)) {
5451 Bindings.addDeclBindingInfo(VD, CBufferOrTBuffer->isCBuffer()
5452 ? ResourceClass::CBuffer
5453 : ResourceClass::SRV);
5454 return;
5455 }
5456
5457 // Unwrap arrays
5458 // FIXME: Calculate array size while unwrapping
5459 const Type *Ty = VD->getType()->getUnqualifiedDesugaredType();
5460 while (Ty->isArrayType()) {
5461 const ArrayType *AT = cast<ArrayType>(Ty);
5463 }
5464
5465 // Resource (or array of resources)
5466 if (const HLSLAttributedResourceType *AttrResType =
5467 HLSLAttributedResourceType::findHandleTypeOnResource(Ty)) {
5468 Bindings.addDeclBindingInfo(VD, AttrResType->getAttrs().ResourceClass);
5469 return;
5470 }
5471
5472 // User defined record type
5473 if (const RecordType *RT = dyn_cast<RecordType>(Ty))
5474 collectResourceBindingsOnUserRecordDecl(VD, RT);
5475}
5476
5477// Walks though the explicit resource binding attributes on the declaration,
5478// and makes sure there is a resource that matched the binding and updates
5479// DeclBindingInfoLists
5480void SemaHLSL::processExplicitBindingsOnDecl(VarDecl *VD) {
5481 assert(VD->hasGlobalStorage() && "expected global variable");
5482
5483 bool HasBinding = false;
5484 for (Attr *A : VD->attrs()) {
5485 if (isa<HLSLVkBindingAttr>(A)) {
5486 HasBinding = true;
5487 if (auto PA = VD->getAttr<HLSLVkPushConstantAttr>())
5488 Diag(PA->getLoc(), diag::err_hlsl_attr_incompatible) << A << PA;
5489 }
5490
5491 HLSLResourceBindingAttr *RBA = dyn_cast<HLSLResourceBindingAttr>(A);
5492 if (!RBA || !RBA->hasRegisterSlot())
5493 continue;
5494 HasBinding = true;
5495
5496 RegisterType RT = RBA->getRegisterType();
5497 assert(RT != RegisterType::I && "invalid or obsolete register type should "
5498 "never have an attribute created");
5499
5500 if (RT == RegisterType::C) {
5501 if (Bindings.hasBindingInfoForDecl(VD))
5502 SemaRef.Diag(VD->getLocation(),
5503 diag::warn_hlsl_user_defined_type_missing_member)
5504 << static_cast<int>(RT);
5505 continue;
5506 }
5507
5508 // Find DeclBindingInfo for this binding and update it, or report error
5509 // if it does not exist (user type does to contain resources with the
5510 // expected resource class).
5512 if (DeclBindingInfo *BI = Bindings.getDeclBindingInfo(VD, RC)) {
5513 // update binding info
5514 BI->setBindingAttribute(RBA, BindingType::Explicit);
5515 } else {
5516 SemaRef.Diag(VD->getLocation(),
5517 diag::warn_hlsl_user_defined_type_missing_member)
5518 << static_cast<int>(RT);
5519 }
5520 }
5521
5522 if (!HasBinding && isResourceRecordTypeOrArrayOf(VD))
5523 SemaRef.Diag(VD->getLocation(), diag::warn_hlsl_implicit_binding);
5524}
5525namespace {
5526class InitListTransformer {
5527 Sema &S;
5528 ASTContext &Ctx;
5529 QualType InitTy;
5530 QualType *DstIt = nullptr;
5531 Expr **ArgIt = nullptr;
5532 // Is wrapping the destination type iterator required? This is only used for
5533 // incomplete array types where we loop over the destination type since we
5534 // don't know the full number of elements from the declaration.
5535 bool Wrap;
5536
5537 bool castInitializer(Expr *E) {
5538 assert(DstIt && "This should always be something!");
5539 if (DstIt == DestTypes.end()) {
5540 if (!Wrap) {
5541 ArgExprs.push_back(E);
5542 // This is odd, but it isn't technically a failure due to conversion, we
5543 // handle mismatched counts of arguments differently.
5544 return true;
5545 }
5546 DstIt = DestTypes.begin();
5547 }
5548 InitializedEntity Entity = InitializedEntity::InitializeParameter(
5549 Ctx, *DstIt, /* Consumed (ObjC) */ false);
5550 ExprResult Res = S.PerformCopyInitialization(Entity, E->getBeginLoc(), E);
5551 if (Res.isInvalid())
5552 return false;
5553 Expr *Init = Res.get();
5554 ArgExprs.push_back(Init);
5555 DstIt++;
5556 return true;
5557 }
5558
5559 bool buildInitializerListImpl(Expr *E) {
5560 // If this is an initialization list, traverse the sub initializers.
5561 if (auto *Init = dyn_cast<InitListExpr>(E)) {
5562 for (auto *SubInit : Init->inits())
5563 if (!buildInitializerListImpl(SubInit))
5564 return false;
5565 return true;
5566 }
5567
5568 // If this is a scalar type, just enqueue the expression.
5569 QualType Ty = E->getType();
5570
5571 if (Ty->isScalarType() || (Ty->isRecordType() && !Ty->isAggregateType()))
5572 return castInitializer(E);
5573
5574 if (auto *VecTy = Ty->getAs<VectorType>()) {
5575 uint64_t Size = VecTy->getNumElements();
5576
5577 QualType SizeTy = Ctx.getSizeType();
5578 uint64_t SizeTySize = Ctx.getTypeSize(SizeTy);
5579 for (uint64_t I = 0; I < Size; ++I) {
5580 auto *Idx = IntegerLiteral::Create(Ctx, llvm::APInt(SizeTySize, I),
5581 SizeTy, SourceLocation());
5582
5584 E, E->getBeginLoc(), Idx, E->getEndLoc());
5585 if (ElExpr.isInvalid())
5586 return false;
5587 if (!castInitializer(ElExpr.get()))
5588 return false;
5589 }
5590 return true;
5591 }
5592 if (auto *MTy = Ty->getAs<ConstantMatrixType>()) {
5593 unsigned Rows = MTy->getNumRows();
5594 unsigned Cols = MTy->getNumColumns();
5595 QualType ElemTy = MTy->getElementType();
5596
5597 for (unsigned R = 0; R < Rows; ++R) {
5598 for (unsigned C = 0; C < Cols; ++C) {
5599 // row index literal
5600 Expr *RowIdx = IntegerLiteral::Create(
5601 Ctx, llvm::APInt(Ctx.getIntWidth(Ctx.IntTy), R), Ctx.IntTy,
5602 E->getBeginLoc());
5603 // column index literal
5604 Expr *ColIdx = IntegerLiteral::Create(
5605 Ctx, llvm::APInt(Ctx.getIntWidth(Ctx.IntTy), C), Ctx.IntTy,
5606 E->getBeginLoc());
5608 E, RowIdx, ColIdx, E->getEndLoc());
5609 if (ElExpr.isInvalid())
5610 return false;
5611 if (!castInitializer(ElExpr.get()))
5612 return false;
5613 ElExpr.get()->setType(ElemTy);
5614 }
5615 }
5616 return true;
5617 }
5618
5619 if (auto *ArrTy = dyn_cast<ConstantArrayType>(Ty.getTypePtr())) {
5620 uint64_t Size = ArrTy->getZExtSize();
5621 QualType SizeTy = Ctx.getSizeType();
5622 uint64_t SizeTySize = Ctx.getTypeSize(SizeTy);
5623 for (uint64_t I = 0; I < Size; ++I) {
5624 auto *Idx = IntegerLiteral::Create(Ctx, llvm::APInt(SizeTySize, I),
5625 SizeTy, SourceLocation());
5627 E, E->getBeginLoc(), Idx, E->getEndLoc());
5628 if (ElExpr.isInvalid())
5629 return false;
5630 if (!buildInitializerListImpl(ElExpr.get()))
5631 return false;
5632 }
5633 return true;
5634 }
5635
5636 if (auto *RD = Ty->getAsCXXRecordDecl()) {
5637 llvm::SmallVector<CXXRecordDecl *> RecordDecls;
5638 RecordDecls.push_back(RD);
5639 while (RecordDecls.back()->getNumBases()) {
5640 CXXRecordDecl *D = RecordDecls.back();
5641 assert(D->getNumBases() == 1 &&
5642 "HLSL doesn't support multiple inheritance");
5643 RecordDecls.push_back(
5645 }
5646 while (!RecordDecls.empty()) {
5647 CXXRecordDecl *RD = RecordDecls.pop_back_val();
5648 for (auto *FD : RD->fields()) {
5649 if (FD->isUnnamedBitField())
5650 continue;
5651 DeclAccessPair Found = DeclAccessPair::make(FD, FD->getAccess());
5652 DeclarationNameInfo NameInfo(FD->getDeclName(), E->getBeginLoc());
5654 E, false, E->getBeginLoc(), CXXScopeSpec(), FD, Found, NameInfo);
5655 if (Res.isInvalid())
5656 return false;
5657 if (!buildInitializerListImpl(Res.get()))
5658 return false;
5659 }
5660 }
5661 }
5662 return true;
5663 }
5664
5665 Expr *generateInitListsImpl(QualType Ty) {
5666 assert(ArgIt != ArgExprs.end() && "Something is off in iteration!");
5667 if (Ty->isScalarType() || (Ty->isRecordType() && !Ty->isAggregateType()))
5668 return *(ArgIt++);
5669
5670 llvm::SmallVector<Expr *> Inits;
5671 Ty = Ty.getDesugaredType(Ctx);
5672 if (Ty->isVectorType() || Ty->isConstantArrayType() ||
5673 Ty->isConstantMatrixType()) {
5674 QualType ElTy;
5675 uint64_t Size = 0;
5676 if (auto *ATy = Ty->getAs<VectorType>()) {
5677 ElTy = ATy->getElementType();
5678 Size = ATy->getNumElements();
5679 } else if (auto *CMTy = Ty->getAs<ConstantMatrixType>()) {
5680 ElTy = CMTy->getElementType();
5681 Size = CMTy->getNumElementsFlattened();
5682 } else {
5683 auto *VTy = cast<ConstantArrayType>(Ty.getTypePtr());
5684 ElTy = VTy->getElementType();
5685 Size = VTy->getZExtSize();
5686 }
5687 for (uint64_t I = 0; I < Size; ++I)
5688 Inits.push_back(generateInitListsImpl(ElTy));
5689 }
5690 if (auto *RD = Ty->getAsCXXRecordDecl()) {
5691 llvm::SmallVector<CXXRecordDecl *> RecordDecls;
5692 RecordDecls.push_back(RD);
5693 while (RecordDecls.back()->getNumBases()) {
5694 CXXRecordDecl *D = RecordDecls.back();
5695 assert(D->getNumBases() == 1 &&
5696 "HLSL doesn't support multiple inheritance");
5697 RecordDecls.push_back(
5699 }
5700 while (!RecordDecls.empty()) {
5701 CXXRecordDecl *RD = RecordDecls.pop_back_val();
5702 for (auto *FD : RD->fields())
5703 if (!FD->isUnnamedBitField())
5704 Inits.push_back(generateInitListsImpl(FD->getType()));
5705 }
5706 }
5707 auto *NewInit = new (Ctx) InitListExpr(Ctx, Inits.front()->getBeginLoc(),
5708 Inits, Inits.back()->getEndLoc());
5709 NewInit->setType(Ty);
5710 return NewInit;
5711 }
5712
5713public:
5714 llvm::SmallVector<QualType, 16> DestTypes;
5715 llvm::SmallVector<Expr *, 16> ArgExprs;
5716 InitListTransformer(Sema &SemaRef, const InitializedEntity &Entity)
5717 : S(SemaRef), Ctx(SemaRef.getASTContext()),
5718 Wrap(Entity.getType()->isIncompleteArrayType()) {
5719 InitTy = Entity.getType().getNonReferenceType();
5720 // When we're generating initializer lists for incomplete array types we
5721 // need to wrap around both when building the initializers and when
5722 // generating the final initializer lists.
5723 if (Wrap) {
5724 assert(InitTy->isIncompleteArrayType());
5725 const IncompleteArrayType *IAT = Ctx.getAsIncompleteArrayType(InitTy);
5726 InitTy = IAT->getElementType();
5727 }
5728 BuildFlattenedTypeList(InitTy, DestTypes);
5729 DstIt = DestTypes.begin();
5730 }
5731
5732 bool buildInitializerList(Expr *E) { return buildInitializerListImpl(E); }
5733
5734 Expr *generateInitLists() {
5735 assert(!ArgExprs.empty() &&
5736 "Call buildInitializerList to generate argument expressions.");
5737 ArgIt = ArgExprs.begin();
5738 if (!Wrap)
5739 return generateInitListsImpl(InitTy);
5740 llvm::SmallVector<Expr *> Inits;
5741 while (ArgIt != ArgExprs.end())
5742 Inits.push_back(generateInitListsImpl(InitTy));
5743
5744 auto *NewInit = new (Ctx) InitListExpr(Ctx, Inits.front()->getBeginLoc(),
5745 Inits, Inits.back()->getEndLoc());
5746 llvm::APInt ArySize(64, Inits.size());
5747 NewInit->setType(Ctx.getConstantArrayType(InitTy, ArySize, nullptr,
5748 ArraySizeModifier::Normal, 0));
5749 return NewInit;
5750 }
5751};
5752} // namespace
5753
5754// Recursively detect any incomplete array anywhere in the type graph,
5755// including arrays, struct fields, and base classes.
5757 Ty = Ty.getCanonicalType();
5758
5759 // Array types
5760 if (const ArrayType *AT = dyn_cast<ArrayType>(Ty)) {
5762 return true;
5764 }
5765
5766 // Record (struct/class) types
5767 if (const auto *RT = Ty->getAs<RecordType>()) {
5768 const RecordDecl *RD = RT->getDecl();
5769
5770 // Walk base classes (for C++ / HLSL structs with inheritance)
5771 if (const auto *CXXRD = dyn_cast<CXXRecordDecl>(RD)) {
5772 for (const CXXBaseSpecifier &Base : CXXRD->bases()) {
5773 if (containsIncompleteArrayType(Base.getType()))
5774 return true;
5775 }
5776 }
5777
5778 // Walk fields
5779 for (const FieldDecl *F : RD->fields()) {
5780 if (containsIncompleteArrayType(F->getType()))
5781 return true;
5782 }
5783 }
5784
5785 return false;
5786}
5787
5789 InitListExpr *Init) {
5790 // If the initializer is a scalar, just return it.
5791 if (Init->getType()->isScalarType())
5792 return true;
5793 ASTContext &Ctx = SemaRef.getASTContext();
5794 InitListTransformer ILT(SemaRef, Entity);
5795
5796 for (unsigned I = 0; I < Init->getNumInits(); ++I) {
5797 Expr *E = Init->getInit(I);
5798 if (E->HasSideEffects(Ctx)) {
5799 QualType Ty = E->getType();
5800 if (Ty->isRecordType())
5801 E = new (Ctx) MaterializeTemporaryExpr(Ty, E, E->isLValue());
5802 E = new (Ctx) OpaqueValueExpr(E->getBeginLoc(), Ty, E->getValueKind(),
5803 E->getObjectKind(), E);
5804 Init->setInit(I, E);
5805 }
5806 if (!ILT.buildInitializerList(E))
5807 return false;
5808 }
5809 size_t ExpectedSize = ILT.DestTypes.size();
5810 size_t ActualSize = ILT.ArgExprs.size();
5811 if (ExpectedSize == 0 && ActualSize == 0)
5812 return true;
5813
5814 // Reject empty initializer if *any* incomplete array exists structurally
5815 if (ActualSize == 0 && containsIncompleteArrayType(Entity.getType())) {
5816 QualType InitTy = Entity.getType().getNonReferenceType();
5817 if (InitTy.hasAddressSpace())
5818 InitTy = SemaRef.getASTContext().removeAddrSpaceQualType(InitTy);
5819
5820 SemaRef.Diag(Init->getBeginLoc(), diag::err_hlsl_incorrect_num_initializers)
5821 << /*TooManyOrFew=*/(int)(ExpectedSize < ActualSize) << InitTy
5822 << /*ExpectedSize=*/ExpectedSize << /*ActualSize=*/ActualSize;
5823 return false;
5824 }
5825
5826 // We infer size after validating legality.
5827 // For incomplete arrays it is completely arbitrary to choose whether we think
5828 // the user intended fewer or more elements. This implementation assumes that
5829 // the user intended more, and errors that there are too few initializers to
5830 // complete the final element.
5831 if (Entity.getType()->isIncompleteArrayType()) {
5832 assert(ExpectedSize > 0 &&
5833 "The expected size of an incomplete array type must be at least 1.");
5834 ExpectedSize =
5835 ((ActualSize + ExpectedSize - 1) / ExpectedSize) * ExpectedSize;
5836 }
5837
5838 // An initializer list might be attempting to initialize a reference or
5839 // rvalue-reference. When checking the initializer we should look through
5840 // the reference.
5841 QualType InitTy = Entity.getType().getNonReferenceType();
5842 if (InitTy.hasAddressSpace())
5843 InitTy = SemaRef.getASTContext().removeAddrSpaceQualType(InitTy);
5844 if (ExpectedSize != ActualSize) {
5845 int TooManyOrFew = ActualSize > ExpectedSize ? 1 : 0;
5846 SemaRef.Diag(Init->getBeginLoc(), diag::err_hlsl_incorrect_num_initializers)
5847 << TooManyOrFew << InitTy << ExpectedSize << ActualSize;
5848 return false;
5849 }
5850
5851 // generateInitListsImpl will always return an InitListExpr here, because the
5852 // scalar case is handled above.
5853 auto *NewInit = cast<InitListExpr>(ILT.generateInitLists());
5854 Init->resizeInits(Ctx, NewInit->getNumInits());
5855 for (unsigned I = 0; I < NewInit->getNumInits(); ++I)
5856 Init->updateInit(Ctx, I, NewInit->getInit(I));
5857 return true;
5858}
5859
5860static QualType ReportMatrixInvalidMember(Sema &S, StringRef Name,
5861 StringRef Expected,
5862 SourceLocation OpLoc,
5863 SourceLocation CompLoc) {
5864 S.Diag(OpLoc, diag::err_builtin_matrix_invalid_member)
5865 << Name << Expected << SourceRange(CompLoc);
5866 return QualType();
5867}
5868
5871 const IdentifierInfo *CompName,
5872 SourceLocation CompLoc) {
5873 const auto *MT = baseType->castAs<ConstantMatrixType>();
5874 StringRef AccessorName = CompName->getName();
5875 assert(!AccessorName.empty() && "Matrix Accessor must have a name");
5876
5877 unsigned Rows = MT->getNumRows();
5878 unsigned Cols = MT->getNumColumns();
5879 bool IsZeroBasedAccessor = false;
5880 unsigned ChunkLen = 0;
5881 if (AccessorName.size() < 2)
5882 return ReportMatrixInvalidMember(S, AccessorName,
5883 "length 4 for zero based: \'_mRC\' or "
5884 "length 3 for one-based: \'_RC\' accessor",
5885 OpLoc, CompLoc);
5886
5887 if (AccessorName[0] == '_') {
5888 if (AccessorName[1] == 'm') {
5889 IsZeroBasedAccessor = true;
5890 ChunkLen = 4; // zero-based: "_mRC"
5891 } else {
5892 ChunkLen = 3; // one-based: "_RC"
5893 }
5894 } else
5896 S, AccessorName, "zero based: \'_mRC\' or one-based: \'_RC\' accessor",
5897 OpLoc, CompLoc);
5898
5899 if (AccessorName.size() % ChunkLen != 0) {
5900 const llvm::StringRef Expected = IsZeroBasedAccessor
5901 ? "zero based: '_mRC' accessor"
5902 : "one-based: '_RC' accessor";
5903
5904 return ReportMatrixInvalidMember(S, AccessorName, Expected, OpLoc, CompLoc);
5905 }
5906
5907 auto isDigit = [](char c) { return c >= '0' && c <= '9'; };
5908 auto isZeroBasedIndex = [](unsigned i) { return i <= 3; };
5909 auto isOneBasedIndex = [](unsigned i) { return i >= 1 && i <= 4; };
5910
5911 bool HasRepeated = false;
5912 SmallVector<bool, 16> Seen(Rows * Cols, false);
5913 unsigned NumComponents = 0;
5914 const char *Begin = AccessorName.data();
5915
5916 for (unsigned I = 0, E = AccessorName.size(); I < E; I += ChunkLen) {
5917 const char *Chunk = Begin + I;
5918 char RowChar = 0, ColChar = 0;
5919 if (IsZeroBasedAccessor) {
5920 // Zero-based: "_mRC"
5921 if (Chunk[0] != '_' || Chunk[1] != 'm') {
5922 char Bad = (Chunk[0] != '_') ? Chunk[0] : Chunk[1];
5924 S, StringRef(&Bad, 1), "\'_m\' prefix",
5925 OpLoc.getLocWithOffset(I + (Bad == Chunk[0] ? 1 : 2)), CompLoc);
5926 }
5927 RowChar = Chunk[2];
5928 ColChar = Chunk[3];
5929 } else {
5930 // One-based: "_RC"
5931 if (Chunk[0] != '_')
5933 S, StringRef(&Chunk[0], 1), "\'_\' prefix",
5934 OpLoc.getLocWithOffset(I + 1), CompLoc);
5935 RowChar = Chunk[1];
5936 ColChar = Chunk[2];
5937 }
5938
5939 // Must be digits.
5940 bool IsDigitsError = false;
5941 if (!isDigit(RowChar)) {
5942 unsigned BadPos = IsZeroBasedAccessor ? 2 : 1;
5943 ReportMatrixInvalidMember(S, StringRef(&RowChar, 1), "row as integer",
5944 OpLoc.getLocWithOffset(I + BadPos + 1),
5945 CompLoc);
5946 IsDigitsError = true;
5947 }
5948
5949 if (!isDigit(ColChar)) {
5950 unsigned BadPos = IsZeroBasedAccessor ? 3 : 2;
5951 ReportMatrixInvalidMember(S, StringRef(&ColChar, 1), "column as integer",
5952 OpLoc.getLocWithOffset(I + BadPos + 1),
5953 CompLoc);
5954 IsDigitsError = true;
5955 }
5956 if (IsDigitsError)
5957 return QualType();
5958
5959 unsigned Row = RowChar - '0';
5960 unsigned Col = ColChar - '0';
5961
5962 bool HasIndexingError = false;
5963 if (IsZeroBasedAccessor) {
5964 // 0-based [0..3]
5965 if (!isZeroBasedIndex(Row)) {
5966 S.Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
5967 << /*row*/ 0 << /*zero-based*/ 0 << SourceRange(CompLoc);
5968 HasIndexingError = true;
5969 }
5970 if (!isZeroBasedIndex(Col)) {
5971 S.Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
5972 << /*col*/ 1 << /*zero-based*/ 0 << SourceRange(CompLoc);
5973 HasIndexingError = true;
5974 }
5975 } else {
5976 // 1-based [1..4]
5977 if (!isOneBasedIndex(Row)) {
5978 S.Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
5979 << /*row*/ 0 << /*one-based*/ 1 << SourceRange(CompLoc);
5980 HasIndexingError = true;
5981 }
5982 if (!isOneBasedIndex(Col)) {
5983 S.Diag(OpLoc, diag::err_hlsl_matrix_element_not_in_bounds)
5984 << /*col*/ 1 << /*one-based*/ 1 << SourceRange(CompLoc);
5985 HasIndexingError = true;
5986 }
5987 // Convert to 0-based after range checking.
5988 --Row;
5989 --Col;
5990 }
5991
5992 if (HasIndexingError)
5993 return QualType();
5994
5995 // Note: matrix swizzle index is hard coded. That means Row and Col can
5996 // potentially be larger than Rows and Cols if matrix size is less than
5997 // the max index size.
5998 bool HasBoundsError = false;
5999 if (Row >= Rows) {
6000 Diag(OpLoc, diag::err_hlsl_matrix_index_out_of_bounds)
6001 << /*Row*/ 0 << Row << Rows << SourceRange(CompLoc);
6002 HasBoundsError = true;
6003 }
6004 if (Col >= Cols) {
6005 Diag(OpLoc, diag::err_hlsl_matrix_index_out_of_bounds)
6006 << /*Col*/ 1 << Col << Cols << SourceRange(CompLoc);
6007 HasBoundsError = true;
6008 }
6009 if (HasBoundsError)
6010 return QualType();
6011
6012 unsigned FlatIndex = Row * Cols + Col;
6013 if (Seen[FlatIndex])
6014 HasRepeated = true;
6015 Seen[FlatIndex] = true;
6016 ++NumComponents;
6017 }
6018 if (NumComponents == 0 || NumComponents > 4) {
6019 S.Diag(OpLoc, diag::err_hlsl_matrix_swizzle_invalid_length)
6020 << NumComponents << SourceRange(CompLoc);
6021 return QualType();
6022 }
6023
6024 QualType ElemTy = MT->getElementType();
6025 if (NumComponents == 1)
6026 return ElemTy;
6027 QualType VT = S.Context.getExtVectorType(ElemTy, NumComponents);
6028 if (HasRepeated)
6029 VK = VK_PRValue;
6030
6031 for (Sema::ExtVectorDeclsType::iterator
6033 E = S.ExtVectorDecls.end();
6034 I != E; ++I) {
6035 if ((*I)->getUnderlyingType() == VT)
6037 /*Qualifier=*/std::nullopt, *I);
6038 }
6039
6040 return VT;
6041}
6042
6044 // If initializing a local resource, track the resource binding it is using
6045 if (VDecl->getType()->isHLSLResourceRecord() && !VDecl->hasGlobalStorage())
6046 trackLocalResource(VDecl, Init);
6047
6048 const HLSLVkConstantIdAttr *ConstIdAttr =
6049 VDecl->getAttr<HLSLVkConstantIdAttr>();
6050 if (!ConstIdAttr)
6051 return true;
6052
6053 ASTContext &Context = SemaRef.getASTContext();
6054
6055 APValue InitValue;
6056 if (!Init->isCXX11ConstantExpr(Context, &InitValue)) {
6057 Diag(VDecl->getLocation(), diag::err_specialization_const);
6058 VDecl->setInvalidDecl();
6059 return false;
6060 }
6061
6062 Builtin::ID BID =
6064
6065 // Argument 1: The ID from the attribute
6066 int ConstantID = ConstIdAttr->getId();
6067 llvm::APInt IDVal(Context.getIntWidth(Context.IntTy), ConstantID);
6068 Expr *IdExpr = IntegerLiteral::Create(Context, IDVal, Context.IntTy,
6069 ConstIdAttr->getLocation());
6070
6071 SmallVector<Expr *, 2> Args = {IdExpr, Init};
6072 Expr *C = SemaRef.BuildBuiltinCallExpr(Init->getExprLoc(), BID, Args);
6073 if (C->getType()->getCanonicalTypeUnqualified() !=
6075 C = SemaRef
6076 .BuildCStyleCastExpr(SourceLocation(),
6077 Context.getTrivialTypeSourceInfo(
6078 Init->getType(), Init->getExprLoc()),
6079 SourceLocation(), C)
6080 .get();
6081 }
6082 Init = C;
6083 return true;
6084}
6085
6087 SourceLocation NameLoc) {
6088 if (!Template)
6089 return QualType();
6090
6091 DeclContext *DC = Template->getDeclContext();
6092 if (!DC->isNamespace() || !cast<NamespaceDecl>(DC)->getIdentifier() ||
6093 cast<NamespaceDecl>(DC)->getName() != "hlsl")
6094 return QualType();
6095
6096 TemplateParameterList *Params = Template->getTemplateParameters();
6097 if (!Params || Params->size() != 1)
6098 return QualType();
6099
6100 if (!Template->isImplicit())
6101 return QualType();
6102
6103 // We manually extract default arguments here instead of letting
6104 // CheckTemplateIdType handle it. This ensures that for resource types that
6105 // lack a default argument (like Buffer), we return a null QualType, which
6106 // triggers the "requires template arguments" error rather than a less
6107 // descriptive "too few template arguments" error.
6108 TemplateArgumentListInfo TemplateArgs(NameLoc, NameLoc);
6109 for (NamedDecl *P : *Params) {
6110 if (auto *TTP = dyn_cast<TemplateTypeParmDecl>(P)) {
6111 if (TTP->hasDefaultArgument()) {
6112 TemplateArgs.addArgument(TTP->getDefaultArgument());
6113 continue;
6114 }
6115 } else if (auto *NTTP = dyn_cast<NonTypeTemplateParmDecl>(P)) {
6116 if (NTTP->hasDefaultArgument()) {
6117 TemplateArgs.addArgument(NTTP->getDefaultArgument());
6118 continue;
6119 }
6120 } else if (auto *TTPD = dyn_cast<TemplateTemplateParmDecl>(P)) {
6121 if (TTPD->hasDefaultArgument()) {
6122 TemplateArgs.addArgument(TTPD->getDefaultArgument());
6123 continue;
6124 }
6125 }
6126 return QualType();
6127 }
6128
6129 return SemaRef.CheckTemplateIdType(
6131 TemplateArgs, nullptr, /*ForNestedNameSpecifier=*/false);
6132}
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....
TokenType getType() const
Returns the token's type, e.g.
FormatToken * Previous
The previous token in the unwrapped line.
Defines the clang::IdentifierInfo, clang::IdentifierTable, and clang::Selector interfaces.
#define X(type, name)
Definition Value.h:97
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)
Definition SemaHLSL.cpp:82
static StringRef createRegisterString(ASTContext &AST, RegisterType RegType, unsigned N)
Definition SemaHLSL.cpp:184
static bool CheckWaveActive(Sema *S, CallExpr *TheCall)
static void createHostLayoutStructForBuffer(Sema &S, HLSLBufferDecl *BufDecl)
Definition SemaHLSL.cpp:609
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)
Definition SemaHLSL.cpp:246
static bool isZeroSizedArray(const ConstantArrayType *CAT)
Definition SemaHLSL.cpp:365
static bool DiagnoseHLSLRegisterAttribute(Sema &S, SourceLocation &ArgLoc, Decl *D, RegisterType RegType, bool SpecifiedSpace)
static bool hasConstantBufferLayout(QualType QT)
static FieldDecl * createFieldForHostLayoutStruct(Sema &S, const Type *Ty, IdentifierInfo *II, CXXRecordDecl *LayoutStruct)
Definition SemaHLSL.cpp:517
static bool CheckUnsignedIntVecRepresentation(Sema *S, SourceLocation Loc, int ArgOrdinal, clang::QualType PassedType)
SampleKind
static bool isInvalidConstantBufferLeafElementType(const Type *Ty)
Definition SemaHLSL.cpp:399
static Builtin::ID getSpecConstBuiltinId(const Type *Type)
Definition SemaHLSL.cpp:150
static bool CheckFloatingOrIntRepresentation(Sema *S, SourceLocation Loc, int ArgOrdinal, clang::QualType PassedType)
static const Type * createHostLayoutType(Sema &S, const Type *Ty)
Definition SemaHLSL.cpp:490
static bool CheckAnyScalarOrVector(Sema *S, CallExpr *TheCall, unsigned ArgIndex)
static const HLSLAttributedResourceType * getResourceArrayHandleType(QualType QT)
Definition SemaHLSL.cpp:381
static IdentifierInfo * getHostLayoutStructName(Sema &S, NamedDecl *BaseDecl, bool MustBeUnique)
Definition SemaHLSL.cpp:455
static void addImplicitBindingAttrToDecl(Sema &S, Decl *D, RegisterType RT, uint32_t ImplicitBindingOrderID)
Definition SemaHLSL.cpp:653
static void SetElementTypeAsReturnType(Sema *S, CallExpr *TheCall, QualType ReturnType)
static unsigned calculateLegacyCbufferSize(const ASTContext &Context, QualType T)
Definition SemaHLSL.cpp:265
static bool CheckLoadLevelBuiltin(Sema &S, CallExpr *TheCall)
static RegisterType getRegisterType(ResourceClass RC)
Definition SemaHLSL.cpp:62
static bool ValidateRegisterNumber(uint64_t SlotNum, Decl *TheDecl, ASTContext &Ctx, RegisterType RegTy)
static bool isVkPipelineBuiltin(const ASTContext &AstContext, FunctionDecl *FD, HLSLAppliedSemanticAttr *Semantic, bool IsInput)
Definition SemaHLSL.cpp:841
static bool CheckVectorElementCount(Sema *S, QualType PassedType, QualType BaseType, unsigned ExpectedCount, SourceLocation Loc)
static bool CheckModifiableLValue(Sema *S, CallExpr *TheCall, unsigned ArgIndex)
static QualType castElement(Sema &S, ExprResult &E, QualType Ty)
static char getRegisterTypeChar(RegisterType RT)
Definition SemaHLSL.cpp:114
static bool CheckNotBoolScalarOrVector(Sema *S, CallExpr *TheCall, unsigned ArgIndex)
static CXXRecordDecl * findRecordDeclInContext(IdentifierInfo *II, DeclContext *DC)
Definition SemaHLSL.cpp:438
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)
Definition SemaHLSL.cpp:132
static CXXRecordDecl * createHostLayoutStruct(Sema &S, CXXRecordDecl *StructDecl)
Definition SemaHLSL.cpp:544
static bool CheckScalarOrVector(Sema *S, CallExpr *TheCall, QualType Scalar, unsigned ArgIndex)
static bool CheckSamplingBuiltin(Sema &S, CallExpr *TheCall, SampleKind Kind)
static bool CheckScalarOrVectorOrMatrix(Sema *S, CallExpr *TheCall, QualType Scalar, unsigned ArgIndex)
static bool CheckFloatRepresentation(Sema *S, SourceLocation Loc, int ArgOrdinal, clang::QualType PassedType)
static bool requiresImplicitBufferLayoutStructure(const CXXRecordDecl *RD)
Definition SemaHLSL.cpp:418
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)
Definition SemaHLSL.cpp:312
static bool IsDefaultBufferConstantDecl(const ASTContext &Ctx, VarDecl *VD)
HLSLResourceBindingAttr::RegisterType RegisterType
Definition SemaHLSL.cpp:57
static CastKind getScalarCastKind(ASTContext &Ctx, QualType DestTy, QualType SrcTy)
static bool CheckGatherBuiltin(Sema &S, CallExpr *TheCall, bool IsCmp)
static bool isValidWaveSizeValue(unsigned Value)
static bool isResourceRecordTypeOrArrayOf(QualType Ty)
Definition SemaHLSL.cpp:372
static bool AccumulateHLSLResourceSlots(QualType Ty, uint64_t &StartSlot, const uint64_t &Limit, const ResourceClass ResClass, ASTContext &Ctx, uint64_t ArrayCount=1)
static bool CheckNoDoubleVectors(Sema *S, SourceLocation Loc, int ArgOrdinal, clang::QualType PassedType)
static bool ValidateMultipleRegisterAnnotations(Sema &S, Decl *TheDecl, RegisterType regType)
static bool CheckTextureSamplerAndLocation(Sema &S, CallExpr *TheCall)
static bool DiagnoseLocalRegisterBinding(Sema &S, SourceLocation &ArgLoc, Decl *D, RegisterType RegType, bool SpecifiedSpace)
This file declares semantic analysis for HLSL constructs.
Defines the clang::SourceLocation class and associated facilities.
Defines various enumerations that describe declaration and type specifiers.
C Language Family Type Representation.
Defines the clang::TypeLoc interface and its subclasses.
C Language Family Type Representation.
static const TypeInfo & getInfo(unsigned id)
Definition Types.cpp:44
__device__ __2f16 float c
return(__x > > __y)|(__x<<(32 - __y))
APValue - This class implements a discriminated union of [uninitialized] [APSInt] [APFloat],...
Definition APValue.h:122
virtual bool HandleTopLevelDecl(DeclGroupRef D)
HandleTopLevelDecl - Handle the specified top-level declaration.
Holds long-lived AST nodes (such as types and decls) that can be referred to throughout the semantic ...
Definition ASTContext.h:226
unsigned getIntWidth(QualType T) const
int getIntegerTypeOrder(QualType LHS, QualType RHS) const
Return the highest ranked integer type, see C99 6.3.1.8p1.
CanQualType FloatTy
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
IdentifierTable & Idents
Definition ASTContext.h:798
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....
CanQualType BoolTy
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.
CanQualType CharTy
CanQualType IntTy
uint64_t getTypeSize(QualType T) const
Return the size of the specified (complete) type T, in bits.
CharUnits getTypeSizeInChars(QualType T) const
Return the size of the specified (complete) type T, in characters.
CanQualType UnsignedIntTy
QualType getTypedefType(ElaboratedTypeKeyword Keyword, NestedNameSpecifier Qualifier, const TypedefNameDecl *Decl, QualType UnderlyingType=QualType(), std::optional< bool > TypeMatchesDeclOrNone=std::nullopt) const
Return the unique reference to the type for the specified typedef-name decl.
llvm::StringRef backupStr(llvm::StringRef S) const
Definition ASTContext.h:880
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
Definition ASTContext.h:917
QualType getHLSLAttributedResourceType(QualType Wrapped, QualType Contained, const HLSLAttributedResourceType::Attributes &Attrs)
QualType getAddrSpaceQualType(QualType T, LangAS AddressSpace) const
Return the uniqued reference to the type for an address space qualified type with the specified type ...
CanQualType getCanonicalTagType(const TagDecl *TD) const
static bool hasSameUnqualifiedType(QualType T1, QualType T2)
Determine whether the given types are equivalent after cvr-qualifiers have been removed.
QualType getConstantMatrixType(QualType ElementType, unsigned NumRows, unsigned NumColumns) const
Return the unique reference to the matrix type of the specified element type and size.
unsigned getTypeAlign(QualType T) const
Return the ABI-specified alignment of a (complete) type T, in bits.
PtrTy get() const
Definition Ownership.h:171
bool isInvalid() const
Definition Ownership.h:167
Represents an array type, per C99 6.7.5.2 - Array Declarators.
Definition TypeBase.h:3772
QualType getElementType() const
Definition TypeBase.h:3784
Attr - This represents one attribute.
Definition Attr.h:46
attr::Kind getKind() const
Definition Attr.h:92
SourceLocation getLocation() const
Definition Attr.h:99
SourceLocation getScopeLoc() const
const IdentifierInfo * getScopeName() const
SourceLocation getLoc() const
const IdentifierInfo * getAttrName() const
Represents a base class of a C++ class.
Definition DeclCXX.h:146
QualType getType() const
Retrieves the type of the base class.
Definition DeclCXX.h:249
Represents a static or instance method of a struct/union/class.
Definition DeclCXX.h:2136
Represents a C++ struct/union/class.
Definition DeclCXX.h:258
bool isHLSLIntangible() const
Returns true if the class contains HLSL intangible type, either as a field or in base class.
Definition DeclCXX.h:1556
static CXXRecordDecl * Create(const ASTContext &C, TagKind TK, DeclContext *DC, SourceLocation StartLoc, SourceLocation IdLoc, IdentifierInfo *Id, CXXRecordDecl *PrevDecl=nullptr)
Definition DeclCXX.cpp:132
void setBases(CXXBaseSpecifier const *const *Bases, unsigned NumBases)
Sets the base classes of this struct or class.
Definition DeclCXX.cpp:184
base_class_iterator bases_end()
Definition DeclCXX.h:617
void completeDefinition() override
Indicates that the definition of this class is now complete.
Definition DeclCXX.cpp:2249
base_class_range bases()
Definition DeclCXX.h:608
unsigned getNumBases() const
Retrieves the number of base classes of this class.
Definition DeclCXX.h:602
base_class_iterator bases_begin()
Definition DeclCXX.h:615
bool isEmpty() const
Determine whether this is an empty class in the sense of (C++11 [meta.unary.prop]).
Definition DeclCXX.h:1186
CallExpr - Represents a function call (C99 6.5.2.2, C++ [expr.call]).
Definition Expr.h:2946
Expr * getArg(unsigned Arg)
getArg - Return the specified argument.
Definition Expr.h:3150
SourceLocation getBeginLoc() const
Definition Expr.h:3280
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.
Definition Expr.cpp:1517
FunctionDecl * getDirectCallee()
If the callee is a FunctionDecl, return it. Otherwise return null.
Definition Expr.h:3129
Expr * getCallee()
Definition Expr.h:3093
unsigned getNumArgs() const
getNumArgs - Return the number of actual arguments to this call.
Definition Expr.h:3137
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.
Definition CharUnits.h:185
Represents the canonical version of C arrays with a specified constant size.
Definition TypeBase.h:3810
bool isZeroSize() const
Return true if the size is zero.
Definition TypeBase.h:3880
llvm::APInt getSize() const
Return the constant array size as an APInt.
Definition TypeBase.h:3866
uint64_t getZExtSize() const
Return the size zero-extended as a uint64_t.
Definition TypeBase.h:3886
Represents a concrete matrix type with constant number of rows and columns.
Definition TypeBase.h:4437
unsigned getNumColumns() const
Returns the number of columns in the matrix.
Definition TypeBase.h:4456
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...
Definition DeclBase.h:1449
bool isNamespace() const
Definition DeclBase.h:2198
lookup_result lookup(DeclarationName Name) const
lookup - Find the declarations (if any) with the given Name in this context.
bool isTranslationUnit() const
Definition DeclBase.h:2185
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.
Definition DeclBase.h:2373
DeclContext * getNonTransparentContext()
A reference to a declared variable, function, enum, etc.
Definition Expr.h:1273
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)
Definition Expr.cpp:488
ValueDecl * getDecl()
Definition Expr.h:1341
Decl - This represents one declaration (or definition), e.g.
Definition DeclBase.h:86
T * getAttr() const
Definition DeclBase.h:573
ASTContext & getASTContext() const LLVM_READONLY
Definition DeclBase.cpp:546
void addAttr(Attr *A)
attr_iterator attr_end() const
Definition DeclBase.h:542
bool isImplicit() const
isImplicit - Indicates whether the declaration was implicitly generated by the implementation.
Definition DeclBase.h:593
void setInvalidDecl(bool Invalid=true)
setInvalidDecl - Indicates the Decl had a semantic error.
Definition DeclBase.cpp:178
bool isInExportDeclContext() const
Whether this declaration was exported in a lexical context.
attr_iterator attr_begin() const
Definition DeclBase.h:539
DeclContext * getNonTransparentDeclContext()
Return the non transparent context.
SourceLocation getLocation() const
Definition DeclBase.h:439
void setImplicit(bool I=true)
Definition DeclBase.h:594
DeclContext * getDeclContext()
Definition DeclBase.h:448
attr_range attrs() const
Definition DeclBase.h:535
AccessSpecifier getAccess() const
Definition DeclBase.h:507
SourceLocation getBeginLoc() const LLVM_READONLY
Definition DeclBase.h:431
void dropAttr()
Definition DeclBase.h:556
bool hasAttr() const
Definition DeclBase.h:577
The name of a declaration.
Represents a ValueDecl that came out of a declarator.
Definition Decl.h:780
SourceLocation getBeginLoc() const LLVM_READONLY
Definition Decl.h:831
This represents one expression.
Definition Expr.h:112
void setType(QualType t)
Definition Expr.h:145
ExprValueKind getValueKind() const
getValueKind - The value kind that this expression produces.
Definition Expr.h:447
Expr * IgnoreParenImpCasts() LLVM_READONLY
Skip past any parentheses and implicit casts which might surround this expression until reaching a fi...
Definition Expr.cpp:3090
Expr * IgnoreParens() LLVM_READONLY
Skip past any parentheses which might surround this expression until reaching a fixed point.
Definition Expr.cpp:3086
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.
Definition Expr.h:284
ExprObjectKind getObjectKind() const
getObjectKind - The object kind that this expression produces.
Definition Expr.h:454
bool HasSideEffects(const ASTContext &Ctx, bool IncludePossibleEffects=true) const
HasSideEffects - This routine returns true for all those expressions which have any effect other than...
Definition Expr.cpp:3670
void setValueKind(ExprValueKind Cat)
setValueKind - Set the value kind produced by this expression.
Definition Expr.h:464
SourceLocation getExprLoc() const LLVM_READONLY
getExprLoc - Return the preferred location for the arrow when diagnosing a problem with a generic exp...
Definition Expr.cpp:277
@ MLV_Valid
Definition Expr.h:306
QualType getType() const
Definition Expr.h:144
ExtVectorType - Extended vector type.
Definition TypeBase.h:4317
Represents difference between two FPOptions values.
Represents a member of a struct/union/class.
Definition Decl.h:3160
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)
Definition Decl.cpp:4701
static FixItHint CreateReplacement(CharSourceRange RemoveRange, StringRef Code)
Create a code modification hint that replaces the given source range with the given code string.
Definition Diagnostic.h:140
Represents a function declaration or definition.
Definition Decl.h:2000
const ParmVarDecl * getParamDecl(unsigned i) const
Definition Decl.h:2797
Stmt * getBody(const FunctionDecl *&Definition) const
Retrieve the body (definition) of the function.
Definition Decl.cpp:3280
bool isThisDeclarationADefinition() const
Returns whether this specific declaration of the function is also a definition that does not contain ...
Definition Decl.h:2314
QualType getReturnType() const
Definition Decl.h:2845
ArrayRef< ParmVarDecl * > parameters() const
Definition Decl.h:2774
bool isTemplateInstantiation() const
Determines if the given function was instantiated from a function template.
Definition Decl.cpp:4258
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.
Definition Decl.cpp:3827
DeclarationNameInfo getNameInfo() const
Definition Decl.h:2211
bool hasBody(const FunctionDecl *&Definition) const
Returns true if the function has a body.
Definition Decl.cpp:3200
bool isDefined(const FunctionDecl *&Definition, bool CheckForPendingFriendDefinition=false) const
Returns true if the function has a definition that does not need to be instantiated.
Definition Decl.cpp:3247
HLSLBufferDecl - Represent a cbuffer or tbuffer declaration.
Definition Decl.h:5196
static HLSLBufferDecl * Create(ASTContext &C, DeclContext *LexicalParent, bool CBuffer, SourceLocation KwLoc, IdentifierInfo *ID, SourceLocation IDLoc, SourceLocation LBrace)
Definition Decl.cpp:5906
void addLayoutStruct(CXXRecordDecl *LS)
Definition Decl.cpp:5946
void setHasValidPackoffset(bool PO)
Definition Decl.h:5241
static HLSLBufferDecl * CreateDefaultCBuffer(ASTContext &C, DeclContext *LexicalParent, ArrayRef< Decl * > DefaultCBufferDecls)
Definition Decl.cpp:5929
buffer_decl_range buffer_decls() const
Definition Decl.h:5271
static HLSLOutArgExpr * Create(const ASTContext &C, QualType Ty, OpaqueValueExpr *Base, OpaqueValueExpr *OpV, Expr *WB, bool IsInOut)
Definition Expr.cpp:5627
static HLSLRootSignatureDecl * Create(ASTContext &C, DeclContext *DC, SourceLocation Loc, IdentifierInfo *ID, llvm::dxbc::RootSignatureVersion Version, ArrayRef< llvm::hlsl::rootsig::RootElement > RootElements)
Definition Decl.cpp:5992
One of these records is kept for each identifier that is lexed.
StringRef getName() const
Return the actual identifier string.
A simple pair of identifier info and location.
SourceLocation getLoc() const
IdentifierInfo * getIdentifierInfo() const
IdentifierInfo & get(StringRef Name)
Return the identifier token info for the specified named identifier.
ImplicitCastExpr - Allows us to explicitly represent implicit type conversions, which have no direct ...
Definition Expr.h:3856
static ImplicitCastExpr * Create(const ASTContext &Context, QualType T, CastKind Kind, Expr *Operand, const CXXCastPath *BasePath, ExprValueKind Cat, FPOptionsOverride FPO)
Definition Expr.cpp:2073
Describes an C or C++ initializer list.
Definition Expr.h:5302
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'.
Definition Expr.cpp:975
iterator begin(Source *source, bool LocalOnly=false)
Represents the results of name lookup.
Definition Lookup.h:147
Represents a prvalue temporary that is written into memory so that a reference can bind to it.
Definition ExprCXX.h:4921
ValueDecl * getMemberDecl() const
Retrieve the member declaration to which this expression refers.
Definition Expr.h:3450
This represents a decl that may have a name.
Definition Decl.h:274
IdentifierInfo * getIdentifier() const
Get the identifier that names this declaration, if there is one.
Definition Decl.h:295
StringRef getName() const
Get the name of identifier for this declaration as a StringRef.
Definition Decl.h:301
DeclarationName getDeclName() const
Get the actual, stored name of the declaration, which may be a special name.
Definition Decl.h:340
A C++ nested-name-specifier augmented with source location information.
OpaqueValueExpr - An expression referring to an opaque object of a fixed type and value class.
Definition Expr.h:1181
Represents a parameter to a function.
Definition Decl.h:1790
ParsedAttr - Represents a syntactic attribute.
Definition ParsedAttr.h:119
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
Definition ParsedAttr.h:389
bool hasParsedType() const
Definition ParsedAttr.h:337
const ParsedType & getTypeArg() const
Definition ParsedAttr.h:459
unsigned getNumArgs() const
getNumArgs - Return the number of actual arguments to this attribute.
Definition ParsedAttr.h:371
bool isArgIdent(unsigned Arg) const
Definition ParsedAttr.h:385
Expr * getArgAsExpr(unsigned Arg) const
Definition ParsedAttr.h:383
AttributeCommonInfo::Kind getKind() const
Definition ParsedAttr.h:610
A (possibly-)qualified type.
Definition TypeBase.h:937
void addRestrict()
Add the restrict qualifier to this QualType.
Definition TypeBase.h:1178
QualType getNonLValueExprType(const ASTContext &Context) const
Determine the type of a (typically non-lvalue) expression with the specified result type.
Definition Type.cpp:3627
QualType getDesugaredType(const ASTContext &Context) const
Return the specified type with any "sugar" removed from the type.
Definition TypeBase.h:1302
bool isNull() const
Return true if this QualType doesn't point to a type yet.
Definition TypeBase.h:1004
const Type * getTypePtr() const
Retrieves a pointer to the underlying (unqualified) type.
Definition TypeBase.h:8431
LangAS getAddressSpace() const
Return the address space of this type.
Definition TypeBase.h:8557
QualType getNonReferenceType() const
If Type is a reference type (e.g., const int&), returns the type that the reference refers to ("const...
Definition TypeBase.h:8616
QualType getCanonicalType() const
Definition TypeBase.h:8483
QualType getUnqualifiedType() const
Retrieve the unqualified variant of the given type, removing as little sugar as possible.
Definition TypeBase.h:8525
bool hasAddressSpace() const
Check if this type has any address space qualifier.
Definition TypeBase.h:8552
Represents a struct/union/class.
Definition Decl.h:4327
field_iterator field_end() const
Definition Decl.h:4533
field_range fields() const
Definition Decl.h:4530
bool field_empty() const
Definition Decl.h:4538
field_iterator field_begin() const
Definition Decl.cpp:5276
bool hasBindingInfoForDecl(const VarDecl *VD) const
Definition SemaHLSL.cpp:220
DeclBindingInfo * getDeclBindingInfo(const VarDecl *VD, ResourceClass ResClass)
Definition SemaHLSL.cpp:206
DeclBindingInfo * addDeclBindingInfo(const VarDecl *VD, ResourceClass ResClass)
Definition SemaHLSL.cpp:193
Scope - A scope is a transient data structure that is used while parsing the program.
Definition Scope.h:41
SemaBase(Sema &S)
Definition SemaBase.cpp:7
ASTContext & getASTContext() const
Definition SemaBase.cpp:9
Sema & SemaRef
Definition SemaBase.h:40
SemaDiagnosticBuilder Diag(SourceLocation Loc, unsigned DiagID)
Emit a diagnostic.
Definition SemaBase.cpp:61
ExprResult ActOnOutParamExpr(ParmVarDecl *Param, Expr *Arg)
HLSLRootSignatureDecl * lookupRootSignatureOverrideDecl(DeclContext *DC) const
bool CanPerformElementwiseCast(Expr *Src, QualType DestType)
void handleWaveSizeAttr(Decl *D, const ParsedAttr &AL)
void handleVkLocationAttr(Decl *D, const ParsedAttr &AL)
HLSLAttributedResourceLocInfo TakeLocForHLSLAttribute(const HLSLAttributedResourceType *RT)
void handleSemanticAttr(Decl *D, const ParsedAttr &AL)
bool CanPerformScalarCast(QualType SrcTy, QualType DestTy)
QualType ProcessResourceTypeAttributes(QualType Wrapped)
void handleShaderAttr(Decl *D, const ParsedAttr &AL)
uint32_t getNextImplicitBindingOrderID()
Definition SemaHLSL.h:228
void CheckEntryPoint(FunctionDecl *FD)
Definition SemaHLSL.cpp:960
void emitLogicalOperatorFixIt(Expr *LHS, Expr *RHS, BinaryOperatorKind Opc)
T * createSemanticAttr(const AttributeCommonInfo &ACI, std::optional< unsigned > Location)
Definition SemaHLSL.h:181
bool initGlobalResourceDecl(VarDecl *VD)
void ActOnEndOfTranslationUnit(TranslationUnitDecl *TU)
bool initGlobalResourceArrayDecl(VarDecl *VD)
HLSLVkConstantIdAttr * mergeVkConstantIdAttr(Decl *D, const AttributeCommonInfo &AL, int Id)
Definition SemaHLSL.cpp:724
HLSLNumThreadsAttr * mergeNumThreadsAttr(Decl *D, const AttributeCommonInfo &AL, int X, int Y, int Z)
Definition SemaHLSL.cpp:690
void deduceAddressSpace(VarDecl *Decl)
std::pair< IdentifierInfo *, bool > ActOnStartRootSignatureDecl(StringRef Signature)
Computes the unique Root Signature identifier from the given signature, then lookup if there is a pre...
void handlePackOffsetAttr(Decl *D, const ParsedAttr &AL)
bool diagnosePositionType(QualType T, const ParsedAttr &AL)
bool handleInitialization(VarDecl *VDecl, Expr *&Init)
bool diagnoseInputIDType(QualType T, const ParsedAttr &AL)
void handleParamModifierAttr(Decl *D, const ParsedAttr &AL)
bool CheckResourceBinOp(BinaryOperatorKind Opc, Expr *LHSExpr, Expr *RHSExpr, SourceLocation Loc)
bool CanPerformAggregateSplatCast(Expr *Src, QualType DestType)
bool IsScalarizedLayoutCompatible(QualType T1, QualType T2) const
QualType ActOnTemplateShorthand(TemplateDecl *Template, SourceLocation NameLoc)
void diagnoseSystemSemanticAttr(Decl *D, const ParsedAttr &AL, std::optional< unsigned > Index)
void handleRootSignatureAttr(Decl *D, const ParsedAttr &AL)
bool CheckCompatibleParameterABI(FunctionDecl *New, FunctionDecl *Old)
QualType handleVectorBinOpConversion(ExprResult &LHS, ExprResult &RHS, QualType LHSType, QualType RHSType, bool IsCompAssign)
QualType checkMatrixComponent(Sema &S, QualType baseType, ExprValueKind &VK, SourceLocation OpLoc, const IdentifierInfo *CompName, SourceLocation CompLoc)
void handleResourceBindingAttr(Decl *D, const ParsedAttr &AL)
bool IsTypedResourceElementCompatible(QualType T1)
bool transformInitList(const InitializedEntity &Entity, InitListExpr *Init)
void handleNumThreadsAttr(Decl *D, const ParsedAttr &AL)
bool ActOnUninitializedVarDecl(VarDecl *D)
void handleVkExtBuiltinInputAttr(Decl *D, const ParsedAttr &AL)
void ActOnTopLevelFunction(FunctionDecl *FD)
Definition SemaHLSL.cpp:793
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)
Definition SemaHLSL.cpp:760
void ActOnFinishBuffer(Decl *Dcl, SourceLocation RBrace)
Definition SemaHLSL.cpp:663
void handleVkBindingAttr(Decl *D, const ParsedAttr &AL)
HLSLParamModifierAttr * mergeParamModifierAttr(Decl *D, const AttributeCommonInfo &AL, HLSLParamModifierAttr::Spelling Spelling)
Definition SemaHLSL.cpp:773
QualType getInoutParameterType(QualType Ty)
SemaHLSL(Sema &S)
Definition SemaHLSL.cpp:224
void handleVkConstantIdAttr(Decl *D, const ParsedAttr &AL)
Decl * ActOnStartBuffer(Scope *BufferScope, bool CBuffer, SourceLocation KwLoc, IdentifierInfo *Ident, SourceLocation IdentLoc, SourceLocation LBrace)
Definition SemaHLSL.cpp:226
HLSLWaveSizeAttr * mergeWaveSizeAttr(Decl *D, const AttributeCommonInfo &AL, int Min, int Max, int Preferred, int SpelledArgsCount)
Definition SemaHLSL.cpp:704
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.
Definition Sema.h:868
@ LookupOrdinaryName
Ordinary name lookup, which finds ordinary names (functions, variables, typedefs, etc....
Definition Sema.h:9385
@ LookupMemberName
Member name lookup, which finds the names of class/struct/union members.
Definition Sema.h:9393
ExtVectorDeclsType ExtVectorDecls
ExtVectorDecls - This is a list all the extended vector types.
Definition Sema.h:4937
ASTContext & Context
Definition Sema.h:1300
ASTContext & getASTContext() const
Definition Sema.h:939
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.
Definition Sema.cpp:757
const LangOptions & getLangOpts() const
Definition Sema.h:932
SemaHLSL & HLSL()
Definition Sema.h:1475
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
Definition Sema.h:942
ASTConsumer & Consumer
Definition Sema.h:1301
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
Definition Stmt.cpp:367
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...
Definition Stmt.cpp:343
SourceLocation getBeginLoc() const LLVM_READONLY
Definition Stmt.cpp:355
StringLiteral - This represents a string literal expression, e.g.
Definition Expr.h:1802
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...
Definition Expr.cpp:1188
void startDefinition()
Starts the definition of this tag declaration.
Definition Decl.cpp:4907
bool isUnion() const
Definition Decl.h:3928
bool isClass() const
Definition Decl.h:3927
Exposes information about the current target.
Definition TargetInfo.h:227
TargetOptions & getTargetOpts() const
Retrieve the target options.
Definition TargetInfo.h:327
const llvm::Triple & getTriple() const
Returns the target triple of the primary target.
StringRef getPlatformName() const
Retrieve the name of the platform as it is used in the availability attribute.
VersionTuple getPlatformMinVersion() const
Retrieve the minimum desired version of the platform, to which the program should be compiled.
std::string HLSLEntry
The entry point name for HLSL shader being compiled as specified by -E.
A convenient class for passing around template argument information.
void addArgument(const TemplateArgumentLoc &Loc)
The base class of all kinds of template declarations (e.g., class, function, etc.).
Stores a list of template parameters for a TemplateDecl and its derived classes.
The top declaration context.
Definition Decl.h:105
SourceLocation getBeginLoc() const
Get the begin source location.
Definition TypeLoc.cpp:193
A container of type source information.
Definition TypeBase.h:8402
TypeLoc getTypeLoc() const
Return the TypeLoc wrapper for the type source info.
Definition TypeLoc.h:267
The base class of the type hierarchy.
Definition TypeBase.h:1866
bool isVoidType() const
Definition TypeBase.h:9034
bool isBooleanType() const
Definition TypeBase.h:9171
bool isIncompleteArrayType() const
Definition TypeBase.h:8775
CXXRecordDecl * getAsCXXRecordDecl() const
Retrieves the CXXRecordDecl that this type refers to, either because the type is a RecordType or beca...
Definition Type.h:26
bool isConstantArrayType() const
Definition TypeBase.h:8771
bool hasIntegerRepresentation() const
Determine whether this type has an integer representation of some sort, e.g., it is an integer type o...
Definition Type.cpp:2083
bool isArrayType() const
Definition TypeBase.h:8767
CXXRecordDecl * castAsCXXRecordDecl() const
Definition Type.h:36
bool isArithmeticType() const
Definition Type.cpp:2374
bool isConstantMatrixType() const
Definition TypeBase.h:8835
bool isHLSLBuiltinIntangibleType() const
Definition TypeBase.h:8979
bool isPointerType() const
Definition TypeBase.h:8668
CanQualType getCanonicalTypeUnqualified() const
bool isIntegerType() const
isIntegerType() does not include complex integers (a GCC extension).
Definition TypeBase.h:9078
const T * castAs() const
Member-template castAs<specific type>.
Definition TypeBase.h:9328
bool isReferenceType() const
Definition TypeBase.h:8692
bool isHLSLIntangibleType() const
Definition Type.cpp:5459
bool isEnumeralType() const
Definition TypeBase.h:8799
bool isScalarType() const
Definition TypeBase.h:9140
bool isIntegralType(const ASTContext &Ctx) const
Determine whether this type is an integral type.
Definition Type.cpp:2120
const Type * getArrayElementTypeNoTypeQual() const
If this is an array type, return the element type of the array, potentially with type qualifiers miss...
Definition Type.cpp:472
QualType getPointeeType() const
If this is a pointer, ObjC object pointer, or block pointer, this returns the respective pointee.
Definition Type.cpp:753
bool hasUnsignedIntegerRepresentation() const
Determine whether this type has an unsigned integer representation of some sort, e....
Definition Type.cpp:2328
bool isAggregateType() const
Determines whether the type is a C++ aggregate type or C aggregate or union type.
Definition Type.cpp:2455
ScalarTypeKind getScalarTypeKind() const
Given that this is a scalar type, classify it.
Definition Type.cpp:2406
bool hasSignedIntegerRepresentation() const
Determine whether this type has an signed integer representation of some sort, e.g....
Definition Type.cpp:2274
bool isMatrixType() const
Definition TypeBase.h:8831
bool isHLSLResourceRecord() const
Definition Type.cpp:5446
bool hasFloatingRepresentation() const
Determine whether this type has a floating-point representation of some sort, e.g....
Definition Type.cpp:2349
bool isVectorType() const
Definition TypeBase.h:8807
bool isRealFloatingType() const
Floating point categories.
Definition Type.cpp:2357
bool isHLSLAttributedResourceType() const
Definition TypeBase.h:8991
@ STK_FloatingComplex
Definition TypeBase.h:2814
@ STK_ObjCObjectPointer
Definition TypeBase.h:2808
@ STK_IntegralComplex
Definition TypeBase.h:2813
@ STK_MemberPointer
Definition TypeBase.h:2809
bool isFloatingType() const
Definition Type.cpp:2341
bool isSamplerT() const
Definition TypeBase.h:8912
const T * getAs() const
Member-template getAs<specific type>'.
Definition TypeBase.h:9261
const Type * getUnqualifiedDesugaredType() const
Return the specified type with any "sugar" removed from the type, removing any typedefs,...
Definition Type.cpp:654
bool isRecordType() const
Definition TypeBase.h:8795
bool isHLSLResourceRecordArray() const
Definition Type.cpp:5450
void setType(QualType newType)
Definition Decl.h:724
QualType getType() const
Definition Decl.h:723
Represents a variable declaration or definition.
Definition Decl.h:926
static VarDecl * Create(ASTContext &C, DeclContext *DC, SourceLocation StartLoc, SourceLocation IdLoc, const IdentifierInfo *Id, QualType T, TypeSourceInfo *TInfo, StorageClass S)
Definition Decl.cpp:2163
void setInitStyle(InitializationStyle Style)
Definition Decl.h:1452
@ CallInit
Call-style initialization (C++98)
Definition Decl.h:934
void setStorageClass(StorageClass SC)
Definition Decl.cpp:2175
bool hasGlobalStorage() const
Returns true for all variables that do not have local storage.
Definition Decl.h:1226
void setInit(Expr *I)
Definition Decl.cpp:2489
StorageClass getStorageClass() const
Returns the storage class as written in the source.
Definition Decl.h:1168
Represents a GCC generic vector type.
Definition TypeBase.h:4225
unsigned getNumElements() const
Definition TypeBase.h:4240
QualType getElementType() const
Definition TypeBase.h:4239
IdentifierInfo * getNameAsIdentifier(ASTContext &AST) const
Defines the clang::TargetInfo interface.
Definition SPIR.cpp:47
uint32_t getResourceDimensions(llvm::dxil::ResourceDimension Dim)
The JSON file list parser is used to communicate input to InstallAPI.
bool isa(CodeGen::Address addr)
Definition Address.h:330
if(T->getSizeExpr()) TRY_TO(TraverseStmt(const_cast< Expr * >(T -> getSizeExpr())))
static bool CheckFloatOrHalfRepresentation(Sema *S, SourceLocation Loc, int ArgOrdinal, clang::QualType PassedType)
Definition SemaSPIRV.cpp:66
@ ICIS_NoInit
No in-class initializer.
Definition Specifiers.h:272
@ TemplateName
The identifier is a template name. FIXME: Add an annotation for that.
Definition Parser.h:61
@ OK_Ordinary
An ordinary object is located at an address in memory.
Definition Specifiers.h:151
static bool CheckAllArgTypesAreCorrect(Sema *S, CallExpr *TheCall, llvm::ArrayRef< llvm::function_ref< bool(Sema *, SourceLocation, int, QualType)> > Checks)
Definition SemaSPIRV.cpp:49
@ AS_public
Definition Specifiers.h:124
@ AS_none
Definition Specifiers.h:127
@ SC_Extern
Definition Specifiers.h:251
@ SC_Static
Definition Specifiers.h:252
@ SC_None
Definition Specifiers.h:250
@ AANT_ArgumentIdentifier
@ Result
The result type of a method or function.
Definition TypeBase.h:905
@ Ordinary
This parameter uses ordinary ABI rules for its type.
Definition Specifiers.h:380
llvm::Expected< QualType > ExpectedType
@ Template
We are parsing a template declaration.
Definition Parser.h:81
LLVM_READONLY bool isDigit(unsigned char c)
Return true if this character is an ASCII digit: [0-9].
Definition CharInfo.h:114
static bool CheckAllArgsHaveSameType(Sema *S, CallExpr *TheCall)
Definition SemaSPIRV.cpp:32
ExprResult ExprError()
Definition Ownership.h:265
@ Type
The name was classified as a type.
Definition Sema.h:564
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.
Definition Specifiers.h:132
@ VK_PRValue
A pr-value expression (in the C++11 taxonomy) produces a temporary value.
Definition Specifiers.h:135
@ VK_LValue
An l-value expression is a reference to an object with independent storage.
Definition Specifiers.h:139
DynamicRecursiveASTVisitorBase< false > DynamicRecursiveASTVisitor
U cast(CodeGen::Address addr)
Definition Address.h:327
@ None
No keyword precedes the qualified type name.
Definition TypeBase.h:5977
ActionResult< Expr * > ExprResult
Definition Ownership.h:249
Visibility
Describes the different kinds of visibility that a declaration may have.
Definition Visibility.h:34
unsigned long uint64_t
unsigned int uint32_t
hash_code hash_value(const clang::dependencies::ModuleID &ID)
__DEVICE__ bool isnan(float __x)
__DEVICE__ _Tp abs(const std::complex< _Tp > &__c)
#define false
Definition stdbool.h:26
Describes how types, statements, expressions, and declarations should be printed.
void setCounterImplicitOrderID(unsigned Value) const
void setImplicitOrderID(unsigned Value) const
const SourceLocation & getLocation() const
Definition SemaHLSL.h:48
const llvm::hlsl::rootsig::RootElement & getElement() const
Definition SemaHLSL.h:47