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