diff options
Diffstat (limited to 'clang/include')
-rw-r--r-- | clang/include/clang/AST/Attr.h | 7 | ||||
-rw-r--r-- | clang/include/clang/AST/InferAlloc.h | 35 | ||||
-rw-r--r-- | clang/include/clang/AST/StmtOpenACC.h | 14 | ||||
-rw-r--r-- | clang/include/clang/Basic/Attr.td | 4 | ||||
-rw-r--r-- | clang/include/clang/Basic/Builtins.td | 7 | ||||
-rw-r--r-- | clang/include/clang/Basic/BuiltinsAMDGPU.def | 9 | ||||
-rw-r--r-- | clang/include/clang/Basic/BuiltinsX86.td | 29 | ||||
-rw-r--r-- | clang/include/clang/Basic/CodeGenOptions.h | 4 | ||||
-rw-r--r-- | clang/include/clang/Basic/DiagnosticASTKinds.td | 6 | ||||
-rw-r--r-- | clang/include/clang/Basic/DiagnosticFrontendKinds.td | 4 | ||||
-rw-r--r-- | clang/include/clang/Basic/DiagnosticSemaKinds.td | 4 | ||||
-rw-r--r-- | clang/include/clang/Basic/LangOptions.h | 8 | ||||
-rw-r--r-- | clang/include/clang/CIR/Dialect/IR/CIROps.td | 37 | ||||
-rw-r--r-- | clang/include/clang/Driver/CommonArgs.h | 3 | ||||
-rw-r--r-- | clang/include/clang/Driver/Options.td | 4 | ||||
-rw-r--r-- | clang/include/clang/Frontend/ASTUnit.h | 5 | ||||
-rw-r--r-- | clang/include/clang/Frontend/CompilerInstance.h | 2 | ||||
-rw-r--r-- | clang/include/clang/Sema/Sema.h | 6 | ||||
-rw-r--r-- | clang/include/clang/Sema/SemaHLSL.h | 30 |
19 files changed, 183 insertions, 35 deletions
diff --git a/clang/include/clang/AST/Attr.h b/clang/include/clang/AST/Attr.h index fe388b9..ce273c1 100644 --- a/clang/include/clang/AST/Attr.h +++ b/clang/include/clang/AST/Attr.h @@ -239,6 +239,8 @@ class HLSLSemanticAttr : public HLSLAnnotationAttr { LLVM_PREFERRED_TYPE(bool) unsigned SemanticExplicitIndex : 1; + Decl *TargetDecl = nullptr; + protected: HLSLSemanticAttr(ASTContext &Context, const AttributeCommonInfo &CommonInfo, attr::Kind AK, bool IsLateParsed, @@ -259,6 +261,11 @@ public: unsigned getSemanticIndex() const { return SemanticIndex; } + bool isSemanticIndexExplicit() const { return SemanticExplicitIndex; } + + void setTargetDecl(Decl *D) { TargetDecl = D; } + Decl *getTargetDecl() const { return TargetDecl; } + // Implement isa/cast/dyncast/etc. static bool classof(const Attr *A) { return A->getKind() >= attr::FirstHLSLSemanticAttr && diff --git a/clang/include/clang/AST/InferAlloc.h b/clang/include/clang/AST/InferAlloc.h new file mode 100644 index 0000000..c3dc302 --- /dev/null +++ b/clang/include/clang/AST/InferAlloc.h @@ -0,0 +1,35 @@ +//===--- InferAlloc.h - Allocation type inference ---------------*- C++ -*-===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// This file defines interfaces for allocation-related type inference. +// +//===----------------------------------------------------------------------===// + +#ifndef LLVM_CLANG_AST_INFERALLOC_H +#define LLVM_CLANG_AST_INFERALLOC_H + +#include "clang/AST/ASTContext.h" +#include "clang/AST/Expr.h" +#include "llvm/Support/AllocToken.h" +#include <optional> + +namespace clang { +namespace infer_alloc { + +/// Infer the possible allocated type from an allocation call expression. +QualType inferPossibleType(const CallExpr *E, const ASTContext &Ctx, + const CastExpr *CastE); + +/// Get the information required for construction of an allocation token ID. +std::optional<llvm::AllocTokenMetadata> +getAllocTokenMetadata(QualType T, const ASTContext &Ctx); + +} // namespace infer_alloc +} // namespace clang + +#endif // LLVM_CLANG_AST_INFERALLOC_H diff --git a/clang/include/clang/AST/StmtOpenACC.h b/clang/include/clang/AST/StmtOpenACC.h index 8b4554e..ae80297 100644 --- a/clang/include/clang/AST/StmtOpenACC.h +++ b/clang/include/clang/AST/StmtOpenACC.h @@ -815,6 +815,20 @@ public: Stmt *getAssociatedStmt() { return OpenACCAssociatedStmtConstruct::getAssociatedStmt(); } + + // A struct to represent a broken-down version of the associated statement, + // providing the information specified in OpenACC3.3 Section 2.12. + struct StmtInfo { + const Expr *V; + const Expr *X; + // Listed as 'expr' in the standard, this is typically a generic expression + // as a component. + const Expr *RefExpr; + // TODO: OpenACC: We should expand this as we're implementing the other + // atomic construct kinds. + }; + + const StmtInfo getAssociatedStmtInfo() const; }; } // namespace clang diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index eb48a0c..749f531 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -787,6 +787,8 @@ class HLSLSemanticAttr<bit Indexable> : HLSLAnnotationAttr { let Spellings = []; let Subjects = SubjectList<[ParmVar, Field, Function]>; let LangOpts = [HLSL]; + let Args = [DeclArgument<Named, "Target">, IntArgument<"SemanticIndex">, + BoolArgument<"SemanticExplicitIndex">]; } /// A target-specific attribute. This class is meant to be used as a mixin @@ -1623,7 +1625,7 @@ def SYCLKernel : InheritableAttr { let Documentation = [SYCLKernelDocs]; } -def DeviceKernel : DeclOrTypeAttr { +def DeviceKernel : InheritableAttr { let Spellings = [Clang<"device_kernel">, Clang<"nvptx_kernel">, Clang<"amdgpu_kernel">, CustomKeyword<"__kernel">, CustomKeyword<"kernel">]; diff --git a/clang/include/clang/Basic/Builtins.td b/clang/include/clang/Basic/Builtins.td index a350acd..a2c2021 100644 --- a/clang/include/clang/Basic/Builtins.td +++ b/clang/include/clang/Basic/Builtins.td @@ -4727,6 +4727,13 @@ def PtrauthStringDiscriminator : Builtin { let Prototype = "size_t(char const*)"; } +// AllocToken builtins. +def InferAllocToken : Builtin { + let Spellings = ["__builtin_infer_alloc_token"]; + let Attributes = [NoThrow, Const, Pure, CustomTypeChecking, Constexpr, UnevaluatedArguments]; + let Prototype = "size_t(...)"; +} + // OpenCL v2.0 s6.13.16, s9.17.3.5 - Pipe functions. // We need the generic prototype, since the packet type could be anything. def ReadPipe : OCLPipeLangBuiltin { diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index 8428fa9..01d121b 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -830,6 +830,15 @@ TARGET_BUILTIN(__builtin_amdgcn_perm_pk16_b4_u4, "V2UiUiUiV2Ui", "nc", "tensor-c TARGET_BUILTIN(__builtin_amdgcn_perm_pk16_b6_u4, "V3UiUiULiV2Ui", "nc", "tensor-cvt-lut-insts") TARGET_BUILTIN(__builtin_amdgcn_perm_pk16_b8_u4, "V4UiULiULiV2Ui", "nc", "tensor-cvt-lut-insts") +TARGET_BUILTIN(__builtin_amdgcn_add_max_i32, "iiiiIb", "nc", "add-min-max-insts") +TARGET_BUILTIN(__builtin_amdgcn_add_max_u32, "UiUiUiUiIb", "nc", "add-min-max-insts") +TARGET_BUILTIN(__builtin_amdgcn_add_min_i32, "iiiiIb", "nc", "add-min-max-insts") +TARGET_BUILTIN(__builtin_amdgcn_add_min_u32, "UiUiUiUiIb", "nc", "add-min-max-insts") +TARGET_BUILTIN(__builtin_amdgcn_pk_add_max_i16, "V2sV2sV2sV2sIb", "nc", "pk-add-min-max-insts") +TARGET_BUILTIN(__builtin_amdgcn_pk_add_max_u16, "V2UsV2UsV2UsV2UsIb", "nc", "pk-add-min-max-insts") +TARGET_BUILTIN(__builtin_amdgcn_pk_add_min_i16, "V2sV2sV2sV2sIb", "nc", "pk-add-min-max-insts") +TARGET_BUILTIN(__builtin_amdgcn_pk_add_min_u16, "V2UsV2UsV2UsV2UsIb", "nc", "pk-add-min-max-insts") + // GFX1250 WMMA builtins TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x4_f32, "V8fIbV2fIbV2fIsV8fIbIb", "nc", "gfx1250-insts,wavefrontsize32") TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x32_bf16, "V8fIbV16yIbV16yIsV8fIbIb", "nc", "gfx1250-insts,wavefrontsize32") diff --git a/clang/include/clang/Basic/BuiltinsX86.td b/clang/include/clang/Basic/BuiltinsX86.td index 54b3ce0..8332eac 100644 --- a/clang/include/clang/Basic/BuiltinsX86.td +++ b/clang/include/clang/Basic/BuiltinsX86.td @@ -123,13 +123,13 @@ let Attributes = [Const, NoThrow, RequiredVectorWidth<128>] in { } } - let Features = "ssse3" in { - def psignb128 : X86Builtin<"_Vector<16, char>(_Vector<16, char>, _Vector<16, char>)">; - def psignw128 : X86Builtin<"_Vector<8, short>(_Vector<8, short>, _Vector<8, short>)">; - def psignd128 : X86Builtin<"_Vector<4, int>(_Vector<4, int>, _Vector<4, int>)">; - } - let Features = "ssse3", Attributes = [NoThrow, Const, Constexpr, RequiredVectorWidth<128>] in { + def psignb128 + : X86Builtin<"_Vector<16, char>(_Vector<16, char>, _Vector<16, char>)">; + def psignw128 + : X86Builtin<"_Vector<8, short>(_Vector<8, short>, _Vector<8, short>)">; + def psignd128 + : X86Builtin<"_Vector<4, int>(_Vector<4, int>, _Vector<4, int>)">; def pmulhrsw128 : X86Builtin<"_Vector<8, short>(_Vector<8, short>, _Vector<8, short>)">; def pmaddubsw128 : X86Builtin<"_Vector<8, short>(_Vector<16, char>, _Vector<16, char>)">; def pshufb128 : X86Builtin<"_Vector<16, char>(_Vector<16, char>, _Vector<16, char>)">; @@ -603,10 +603,9 @@ let Features = "avx2", Attributes = [NoThrow, Const, RequiredVectorWidth<256>] i def mpsadbw256 : X86Builtin<"_Vector<32, char>(_Vector<32, char>, _Vector<32, char>, _Constant char)">; def palignr256 : X86Builtin<"_Vector<32, char>(_Vector<32, char>, " "_Vector<32, char>, _Constant int)">; - def psadbw256 : X86Builtin<"_Vector<4, long long int>(_Vector<32, char>, _Vector<32, char>)">; - def psignb256 : X86Builtin<"_Vector<32, char>(_Vector<32, char>, _Vector<32, char>)">; - def psignw256 : X86Builtin<"_Vector<16, short>(_Vector<16, short>, _Vector<16, short>)">; - def psignd256 : X86Builtin<"_Vector<8, int>(_Vector<8, int>, _Vector<8, int>)">; + def psadbw256 + : X86Builtin< + "_Vector<4, long long int>(_Vector<32, char>, _Vector<32, char>)">; def psllw256 : X86Builtin<"_Vector<16, short>(_Vector<16, short>, _Vector<8, short>)">; def pslldqi256_byteshift : X86Builtin<"_Vector<32, char>(_Vector<32, char>, _Constant int)">; def pslld256 : X86Builtin<"_Vector<8, int>(_Vector<8, int>, _Vector<4, int>)">; @@ -677,7 +676,15 @@ let Features = "avx2", Attributes = [NoThrow, Const, Constexpr, RequiredVectorWi def phsubw256 : X86Builtin<"_Vector<16, short>(_Vector<16, short>, _Vector<16, short>)">; def phsubd256 : X86Builtin<"_Vector<8, int>(_Vector<8, int>, _Vector<8, int>)">; def phsubsw256 : X86Builtin<"_Vector<16, short>(_Vector<16, short>, _Vector<16, short>)">; - + + def psignb256 + : X86Builtin<"_Vector<32, char>(_Vector<32, char>, _Vector<32, char>)">; + def psignw256 + : X86Builtin< + "_Vector<16, short>(_Vector<16, short>, _Vector<16, short>)">; + def psignd256 + : X86Builtin<"_Vector<8, int>(_Vector<8, int>, _Vector<8, int>)">; + def pshuflw256 : X86Builtin<"_Vector<16, short>(_Vector<16, short>, _Constant int)">; def pshufhw256 : X86Builtin<"_Vector<16, short>(_Vector<16, short>, _Constant int)">; def pshufd256 : X86Builtin<"_Vector<8, int>(_Vector<8, int>, _Constant int)">; diff --git a/clang/include/clang/Basic/CodeGenOptions.h b/clang/include/clang/Basic/CodeGenOptions.h index cae06c3..5d5cf25 100644 --- a/clang/include/clang/Basic/CodeGenOptions.h +++ b/clang/include/clang/Basic/CodeGenOptions.h @@ -447,10 +447,6 @@ public: std::optional<double> AllowRuntimeCheckSkipHotCutoff; - /// Maximum number of allocation tokens (0 = no max), nullopt if none set (use - /// pass default). - std::optional<uint64_t> AllocTokenMax; - /// List of backend command-line options for -fembed-bitcode. std::vector<uint8_t> CmdArgs; diff --git a/clang/include/clang/Basic/DiagnosticASTKinds.td b/clang/include/clang/Basic/DiagnosticASTKinds.td index 0be9146f..5c462f9 100644 --- a/clang/include/clang/Basic/DiagnosticASTKinds.td +++ b/clang/include/clang/Basic/DiagnosticASTKinds.td @@ -403,6 +403,12 @@ def note_constexpr_assumption_failed : Note< def note_constexpr_countzeroes_zero : Note< "evaluation of %select{__builtin_elementwise_clzg|__builtin_elementwise_ctzg}0 " "with a zero value is undefined">; +def note_constexpr_infer_alloc_token_type_inference_failed : Note< + "could not infer allocation type for __builtin_infer_alloc_token">; +def note_constexpr_infer_alloc_token_no_metadata : Note< + "could not get token metadata for inferred type">; +def note_constexpr_infer_alloc_token_stateful_mode : Note< + "stateful alloc token mode not supported in constexpr">; def err_experimental_clang_interp_failed : Error< "the experimental clang interpreter failed to evaluate an expression">; diff --git a/clang/include/clang/Basic/DiagnosticFrontendKinds.td b/clang/include/clang/Basic/DiagnosticFrontendKinds.td index 64391de..9e34416 100644 --- a/clang/include/clang/Basic/DiagnosticFrontendKinds.td +++ b/clang/include/clang/Basic/DiagnosticFrontendKinds.td @@ -404,10 +404,6 @@ def warn_hlsl_langstd_minimal : "recommend using %1 instead">, InGroup<HLSLDXCCompat>; -def err_hlsl_semantic_missing : Error<"semantic annotations must be present " - "for all input and outputs of an entry " - "function or patch constant function">; - // ClangIR frontend errors def err_cir_to_cir_transform_failed : Error< "CIR-to-CIR transformation failed">, DefaultFatal; diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 5ff4cc4..13f0d59 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -4126,6 +4126,9 @@ def warn_missing_sdksettings_for_availability_checking : Warning< "%0 availability is ignored without a valid 'SDKSettings.json' in the SDK">, InGroup<DiagGroup<"ignored-availability-without-sdk-settings">>; +def err_hidden_device_kernel + : Error<"%0 is specified as a device kernel but it is not externally visible">; + // Thread Safety Attributes def warn_thread_attribute_ignored : Warning< "ignoring %0 attribute because its argument is invalid">, @@ -13170,6 +13173,7 @@ def err_hlsl_duplicate_parameter_modifier : Error<"duplicate parameter modifier def err_hlsl_missing_semantic_annotation : Error< "semantic annotations must be present for all parameters of an entry " "function or patch constant function">; +def note_hlsl_semantic_used_here : Note<"%0 used here">; def err_hlsl_unknown_semantic : Error<"unknown HLSL semantic %0">; def err_hlsl_semantic_output_not_supported : Error<"semantic %0 does not support output">; diff --git a/clang/include/clang/Basic/LangOptions.h b/clang/include/clang/Basic/LangOptions.h index 260a753..8aa89d8 100644 --- a/clang/include/clang/Basic/LangOptions.h +++ b/clang/include/clang/Basic/LangOptions.h @@ -25,6 +25,7 @@ #include "llvm/ADT/FloatingPointMode.h" #include "llvm/ADT/StringRef.h" #include "llvm/BinaryFormat/DXContainer.h" +#include "llvm/Support/AllocToken.h" #include "llvm/TargetParser/Triple.h" #include <optional> #include <string> @@ -565,6 +566,13 @@ public: bool AtomicFineGrainedMemory = false; bool AtomicIgnoreDenormalMode = false; + /// Maximum number of allocation tokens (0 = no max), nullopt if none set (use + /// target default). + std::optional<uint64_t> AllocTokenMax; + + /// The allocation token mode. + std::optional<llvm::AllocTokenMode> AllocTokenMode; + LangOptions(); /// Set language defaults for the given input language and diff --git a/clang/include/clang/CIR/Dialect/IR/CIROps.td b/clang/include/clang/CIR/Dialect/IR/CIROps.td index 86d09d7..2b361ed 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIROps.td +++ b/clang/include/clang/CIR/Dialect/IR/CIROps.td @@ -4053,6 +4053,43 @@ def CIR_ExpectOp : CIR_Op<"expect", [ } //===----------------------------------------------------------------------===// +// PrefetchOp +//===----------------------------------------------------------------------===// + +def CIR_PrefetchOp : CIR_Op<"prefetch"> { + let summary = "Prefetch operation"; + let description = [{ + The `cir.prefetch` operation is a hint to the code generator to insert a + prefetch instruction if supported; otherwise, it is a noop. Prefetches + have no effect on the behavior of the program but can change its + performance characteristics. + + ```mlir + cir.prefetch(%0 : !cir.ptr<!void>) locality(1) write + ``` + + $locality is a temporal locality specifier ranging from (0) - no locality, + to (3) - extremely local, keep in cache. If $locality is not present, the + default value is 3. + + $isWrite specifies whether the prefetch is for a 'read' or 'write'. If + $isWrite is not specified, it means that prefetch is prepared for 'read'. + }]; + + let arguments = (ins CIR_VoidPtrType:$addr, + DefaultValuedAttr<ConfinedAttr<I32Attr, [IntMinValue<0>, IntMaxValue<3>]>, + "3">:$locality, + UnitAttr:$isWrite); + + let assemblyFormat = [{ + (`write` $isWrite^) : (`read`)? + `locality` `(` $locality `)` + $addr `:` qualified(type($addr)) + attr-dict + }]; +} + +//===----------------------------------------------------------------------===// // PtrDiffOp //===----------------------------------------------------------------------===// diff --git a/clang/include/clang/Driver/CommonArgs.h b/clang/include/clang/Driver/CommonArgs.h index 23426c0..ac17d62 100644 --- a/clang/include/clang/Driver/CommonArgs.h +++ b/clang/include/clang/Driver/CommonArgs.h @@ -76,6 +76,9 @@ void SplitDebugInfo(const ToolChain &TC, Compilation &C, const Tool &T, const JobAction &JA, const llvm::opt::ArgList &Args, const InputInfo &Output, const char *OutFile); +void addDTLTOOptions(const ToolChain &ToolChain, const llvm::opt::ArgList &Args, + llvm::opt::ArgStringList &CmdArgs); + void addLTOOptions(const ToolChain &ToolChain, const llvm::opt::ArgList &Args, llvm::opt::ArgStringList &CmdArgs, const InputInfo &Output, const InputInfoList &Inputs, bool IsThinLTO); diff --git a/clang/include/clang/Driver/Options.td b/clang/include/clang/Driver/Options.td index 7ae153d..0c9584f 100644 --- a/clang/include/clang/Driver/Options.td +++ b/clang/include/clang/Driver/Options.td @@ -2751,6 +2751,10 @@ def falloc_token_max_EQ : Joined<["-"], "falloc-token-max=">, MetaVarName<"<N>">, HelpText<"Limit to maximum N allocation tokens (0 = no max)">; +def falloc_token_mode_EQ : Joined<["-"], "falloc-token-mode=">, + Group<f_Group>, Visibility<[CC1Option]>, + HelpText<"Set the allocation token mode (experimental)">; + def fallow_runtime_check_skip_hot_cutoff_EQ : Joined<["-"], "fallow-runtime-check-skip-hot-cutoff=">, Group<f_clang_Group>, diff --git a/clang/include/clang/Frontend/ASTUnit.h b/clang/include/clang/Frontend/ASTUnit.h index f66df89..3cea159 100644 --- a/clang/include/clang/Frontend/ASTUnit.h +++ b/clang/include/clang/Frontend/ASTUnit.h @@ -499,6 +499,11 @@ public: return *PPOpts; } + IntrusiveRefCntPtr<llvm::vfs::FileSystem> getVirtualFileSystemPtr() { + // FIXME: Don't defer VFS ownership to the FileManager. + return FileMgr->getVirtualFileSystemPtr(); + } + const FileManager &getFileManager() const { return *FileMgr; } FileManager &getFileManager() { return *FileMgr; } IntrusiveRefCntPtr<FileManager> getFileManagerPtr() { return FileMgr; } diff --git a/clang/include/clang/Frontend/CompilerInstance.h b/clang/include/clang/Frontend/CompilerInstance.h index 44fff69..2403cbb 100644 --- a/clang/include/clang/Frontend/CompilerInstance.h +++ b/clang/include/clang/Frontend/CompilerInstance.h @@ -460,7 +460,7 @@ public: FileMgr.resetWithoutRelease(); } - /// Replace the current file manager and virtual file system. + /// Replace the current file manager. void setFileManager(IntrusiveRefCntPtr<FileManager> Value); /// @} diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 87b96c2..189798f 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -1321,15 +1321,11 @@ public: /// Callback to the parser to parse templated functions when needed. typedef void LateTemplateParserCB(void *P, LateParsedTemplate &LPT); - typedef void LateTemplateParserCleanupCB(void *P); LateTemplateParserCB *LateTemplateParser; - LateTemplateParserCleanupCB *LateTemplateParserCleanup; void *OpaqueParser; - void SetLateTemplateParser(LateTemplateParserCB *LTP, - LateTemplateParserCleanupCB *LTPCleanup, void *P) { + void SetLateTemplateParser(LateTemplateParserCB *LTP, void *P) { LateTemplateParser = LTP; - LateTemplateParserCleanup = LTPCleanup; OpaqueParser = P; } diff --git a/clang/include/clang/Sema/SemaHLSL.h b/clang/include/clang/Sema/SemaHLSL.h index f9d3a4ea..8c3b6ae 100644 --- a/clang/include/clang/Sema/SemaHLSL.h +++ b/clang/include/clang/Sema/SemaHLSL.h @@ -130,9 +130,6 @@ public: bool ActOnUninitializedVarDecl(VarDecl *D); void ActOnEndOfTranslationUnit(TranslationUnitDecl *TU); void CheckEntryPoint(FunctionDecl *FD); - bool isSemanticValid(FunctionDecl *FD, DeclaratorDecl *D); - void CheckSemanticAnnotation(FunctionDecl *EntryPoint, const Decl *Param, - const HLSLAnnotationAttr *AnnotationAttr); bool CheckResourceBinOp(BinaryOperatorKind Opc, Expr *LHSExpr, Expr *RHSExpr, SourceLocation Loc); void DiagnoseAttrStageMismatch( @@ -179,17 +176,17 @@ public: bool handleResourceTypeAttr(QualType T, const ParsedAttr &AL); template <typename T> - T *createSemanticAttr(const ParsedAttr &AL, + T *createSemanticAttr(const AttributeCommonInfo &ACI, NamedDecl *TargetDecl, std::optional<unsigned> Location) { - T *Attr = ::new (getASTContext()) T(getASTContext(), AL); - if (Attr->isSemanticIndexable()) - Attr->setSemanticIndex(Location ? *Location : 0); - else if (Location.has_value()) { + T *Attr = + ::new (getASTContext()) T(getASTContext(), ACI, TargetDecl, + Location.value_or(0), Location.has_value()); + + if (!Attr->isSemanticIndexable() && Location.has_value()) { Diag(Attr->getLocation(), diag::err_hlsl_semantic_indexing_not_supported) << Attr->getAttrName()->getName(); return nullptr; } - return Attr; } @@ -247,10 +244,25 @@ private: IdentifierInfo *RootSigOverrideIdent = nullptr; + struct SemanticInfo { + HLSLSemanticAttr *Semantic; + std::optional<uint32_t> Index; + }; + private: void collectResourceBindingsOnVarDecl(VarDecl *D); void collectResourceBindingsOnUserRecordDecl(const VarDecl *VD, const RecordType *RT); + + void checkSemanticAnnotation(FunctionDecl *EntryPoint, const Decl *Param, + const HLSLSemanticAttr *SemanticAttr); + HLSLSemanticAttr *createSemantic(const SemanticInfo &Semantic, + DeclaratorDecl *TargetDecl); + bool determineActiveSemanticOnScalar(FunctionDecl *FD, DeclaratorDecl *D, + SemanticInfo &ActiveSemantic); + bool determineActiveSemantic(FunctionDecl *FD, DeclaratorDecl *D, + SemanticInfo &ActiveSemantic); + void processExplicitBindingsOnDecl(VarDecl *D); void diagnoseAvailabilityViolations(TranslationUnitDecl *TU); |