aboutsummaryrefslogtreecommitdiff
path: root/clang/include
diff options
context:
space:
mode:
Diffstat (limited to 'clang/include')
-rw-r--r--clang/include/clang/AST/Attr.h7
-rw-r--r--clang/include/clang/AST/InferAlloc.h35
-rw-r--r--clang/include/clang/AST/StmtOpenACC.h14
-rw-r--r--clang/include/clang/Basic/Attr.td4
-rw-r--r--clang/include/clang/Basic/Builtins.td7
-rw-r--r--clang/include/clang/Basic/BuiltinsAMDGPU.def9
-rw-r--r--clang/include/clang/Basic/BuiltinsX86.td29
-rw-r--r--clang/include/clang/Basic/CodeGenOptions.h4
-rw-r--r--clang/include/clang/Basic/DiagnosticASTKinds.td6
-rw-r--r--clang/include/clang/Basic/DiagnosticFrontendKinds.td4
-rw-r--r--clang/include/clang/Basic/DiagnosticSemaKinds.td4
-rw-r--r--clang/include/clang/Basic/LangOptions.h8
-rw-r--r--clang/include/clang/CIR/Dialect/IR/CIROps.td37
-rw-r--r--clang/include/clang/Driver/CommonArgs.h3
-rw-r--r--clang/include/clang/Driver/Options.td4
-rw-r--r--clang/include/clang/Frontend/ASTUnit.h5
-rw-r--r--clang/include/clang/Frontend/CompilerInstance.h2
-rw-r--r--clang/include/clang/Sema/Sema.h6
-rw-r--r--clang/include/clang/Sema/SemaHLSL.h30
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);