https://github.com/Pierre-vh updated https://github.com/llvm/llvm-project/pull/207687
>From 78e32b692c5cbb4b6a20abf2970e378fe47a5e48 Mon Sep 17 00:00:00 2001 From: pvanhout <[email protected]> Date: Wed, 15 Jul 2026 13:04:14 +0200 Subject: [PATCH] [clang][AMDGPU] Clean-up handling of named barrier type - Allow the type in struct/classes in very limited circumstances. The goal is to enable creating trivial wrappers around the named barrier variable, but ensure we can't get into situations where things would get awkward. Currently this means we only allow the named barrier in RecordDecls with exactly 1 field, that have no base class, and are not inherited. - Use a `amdgpu_barrier` LangAS for this type that currently maps to the local AS. This allows easy switching to the barrier AS in a future patch. --- clang/include/clang/AST/TypeBase.h | 6 ++ clang/include/clang/Basic/AddressSpaces.h | 3 + clang/include/clang/Basic/Attr.td | 9 +- .../clang/Basic/DiagnosticSemaKinds.td | 13 +++ clang/include/clang/Basic/TargetInfo.h | 7 ++ clang/include/clang/Sema/SemaAMDGPU.h | 3 + clang/lib/AST/ASTContext.cpp | 8 +- clang/lib/AST/Type.cpp | 27 +++++- clang/lib/AST/TypePrinter.cpp | 2 + clang/lib/Basic/TargetInfo.cpp | 3 + clang/lib/Basic/Targets/AArch64.h | 1 + clang/lib/Basic/Targets/AMDGPU.cpp | 3 + clang/lib/Basic/Targets/DirectX.h | 1 + clang/lib/Basic/Targets/NVPTX.h | 1 + clang/lib/Basic/Targets/SPIR.h | 3 + clang/lib/Basic/Targets/SystemZ.h | 3 +- clang/lib/Basic/Targets/TCE.h | 1 + clang/lib/Basic/Targets/WebAssembly.h | 1 + clang/lib/Basic/Targets/X86.h | 1 + clang/lib/CodeGen/CodeGenModule.cpp | 3 + clang/lib/Sema/Sema.cpp | 9 +- clang/lib/Sema/SemaAMDGPU.cpp | 85 +++++++++++++++++++ clang/lib/Sema/SemaDecl.cpp | 4 + clang/test/CodeGenHIP/amdgpu-barrier-type.hip | 42 +++++---- clang/test/SemaCXX/amdgpu-barrier.cpp | 71 ++++++++++++++++ clang/test/SemaHIP/amdgpu-barrier.hip | 72 ++++++++++++++++ clang/test/SemaOpenCL/amdgpu-barrier.cl | 24 ++++++ .../SemaTemplate/address_space-dependent.cpp | 4 +- 28 files changed, 376 insertions(+), 34 deletions(-) diff --git a/clang/include/clang/AST/TypeBase.h b/clang/include/clang/AST/TypeBase.h index c9658775f0470..2182cd79144fa 100644 --- a/clang/include/clang/AST/TypeBase.h +++ b/clang/include/clang/AST/TypeBase.h @@ -2813,6 +2813,12 @@ class alignas(TypeAlignment) Type : public ExtQualsTypeCommonBase { /// Check if the type is the CUDA device builtin texture type. bool isCUDADeviceBuiltinTextureType() const; + /// Check if the type is the AMDGPU named barrier type, or an array thereof. + bool isAMDGPUNamedBarrierType() const; + /// Check if the type is the AMDGPU named barrier type/a RecordType of a named + /// barrier wrapper, or an array thereof. + bool isAMDGPUNamedBarrierTypeOrWrapper() const; + /// Return the implicit lifetime for this type, which must not be dependent. Qualifiers::ObjCLifetime getObjCARCImplicitLifetime() const; diff --git a/clang/include/clang/Basic/AddressSpaces.h b/clang/include/clang/Basic/AddressSpaces.h index a941805423bca..339fc493bfbaf 100644 --- a/clang/include/clang/Basic/AddressSpaces.h +++ b/clang/include/clang/Basic/AddressSpaces.h @@ -68,6 +68,9 @@ enum class LangAS : unsigned { // Wasm specific address spaces. wasm_funcref, + // HIP-specific address spaces + amdgpu_barrier, + // This denotes the count of language-specific address spaces and also // the offset added to the target-specific address spaces, which are usually // specified by address space attributes __attribute__(address_space(n))). diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 1a5bd2301dfc8..f486bd176c9f8 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -2532,6 +2532,13 @@ def AMDGPUMaxNumWorkGroups : InheritableAttr { let Subjects = SubjectList<[Function], ErrorDiag, "kernel functions">; } +def AMDGPUNamedBarrierWrapper : InheritableAttr { + let Spellings = []; + let Args = []; + let Documentation = [InternalOnly]; + let SemaHandler = 0; +} + def BPFPreserveAccessIndex : InheritableAttr, TargetSpecificAttr<TargetBPF> { let Spellings = [Clang<"preserve_access_index">]; @@ -5316,7 +5323,7 @@ def HLSLVkLocation : HLSLAnnotationAttr { } // `row_major` / `column_major` are HLSL keywords that select the in-memory -// layout of a matrix-typed declaration. +// layout of a matrix-typed declaration. def HLSLRowMajor : TypeAttr { let Spellings = [CustomKeyword<"row_major">]; let LangOpts = [HLSL]; diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 3ff7e30d9f1f7..effa7516f9a5a 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -14295,6 +14295,19 @@ def note_acc_reduction_combiner_forming : Note<"while forming %select{|binary operator '%1'|conditional " "operator|final assignment operator}0">; +// AMDGCN type diagnostics +def err_amdgcn_invalid_field_not_a_wrapper : Error< + "fields of type %0 are only allowed in named barrier wrappers">; +def note_amdgcn_not_a_named_barrier_wrapper_too_many_fields : Note< + "%0 is not a named barrier wrapper because it has more than one field">; + +def err_amdgcn_named_barrier_wrapper_non_standard_layout : Error< + "named barrier wrapper %0 must have a C++11 standard layout">; +def note_amdgcn_named_barrier_reason_field : Note< + "%0 is a named barrier wrapper because it has a named barrier or named barrier wrapper field %1 found here">; +def note_amdgcn_named_barrier_reason_inherited : Note< + "%0 is a named barrier wrapper because it inherits from named barrier wrapper %1">; + // AMDGCN builtins diagnostics def err_amdgcn_load_lds_size_invalid_value : Error<"invalid size value">; def note_amdgcn_load_lds_size_valid_value : Note<"size must be %select{1, 2, or 4|1, 2, 4, 12 or 16}0">; diff --git a/clang/include/clang/Basic/TargetInfo.h b/clang/include/clang/Basic/TargetInfo.h index 3aba4d261a651..871666af1da77 100644 --- a/clang/include/clang/Basic/TargetInfo.h +++ b/clang/include/clang/Basic/TargetInfo.h @@ -287,6 +287,9 @@ class TargetInfo : public TransferrableTargetInfo, LLVM_PREFERRED_TYPE(bool) unsigned HasUnalignedAccess : 1; + LLVM_PREFERRED_TYPE(bool) + unsigned HasAMDGPUTypes : 1; + unsigned ARMCDECoprocMask : 8; unsigned MaxOpenCLWorkGroupSize; @@ -1073,6 +1076,10 @@ class TargetInfo : public TransferrableTargetInfo, /// available on this target. bool hasAArch64ACLETypes() const { return HasAArch64ACLETypes; } + /// Returns whether or not the AMDGPU built-in types are + /// available on this target. + bool hasAMDGPUTypes() const { return HasAMDGPUTypes; } + /// Returns whether or not the RISC-V V built-in types are /// available on this target. bool hasRISCVVTypes() const { return HasRISCVVTypes; } diff --git a/clang/include/clang/Sema/SemaAMDGPU.h b/clang/include/clang/Sema/SemaAMDGPU.h index a6205534e0de3..897b1a03dc10b 100644 --- a/clang/include/clang/Sema/SemaAMDGPU.h +++ b/clang/include/clang/Sema/SemaAMDGPU.h @@ -89,6 +89,9 @@ class SemaAMDGPU : public SemaBase { void AddPotentiallyUnguardedBuiltinUser(FunctionDecl *FD); bool HasPotentiallyUnguardedBuiltinUsage(FunctionDecl *FD) const; void DiagnoseUnguardedBuiltinUsage(FunctionDecl *FD); + + /// Called in `ActOnFields` - whenever a C/C++ Record is being finalized. + void checkNamedBarrierWrapper(RecordDecl *R); }; } // namespace clang diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp index 2228811546c0f..6bc5e7a607e5b 100644 --- a/clang/lib/AST/ASTContext.cpp +++ b/clang/lib/AST/ASTContext.cpp @@ -1478,13 +1478,7 @@ void ASTContext::InitBuiltinTypes(const TargetInfo &Target, #include "clang/Basic/WebAssemblyReferenceTypes.def" } - if (Target.getTriple().isAMDGPU() || - (Target.getTriple().isSPIRV() && - Target.getTriple().getVendor() == llvm::Triple::AMD) || - (AuxTarget && - (AuxTarget->getTriple().isAMDGPU() || - ((AuxTarget->getTriple().isSPIRV() && - AuxTarget->getTriple().getVendor() == llvm::Triple::AMD))))) { + if (Target.hasAMDGPUTypes() || (AuxTarget && (AuxTarget->hasAMDGPUTypes()))) { #define AMDGPU_TYPE(Name, Id, SingletonId, Width, Align) \ InitBuiltinType(SingletonId, BuiltinType::Id); #include "clang/Basic/AMDGPUTypes.def" diff --git a/clang/lib/AST/Type.cpp b/clang/lib/AST/Type.cpp index 42d148715bc40..e51e7de9f176a 100644 --- a/clang/lib/AST/Type.cpp +++ b/clang/lib/AST/Type.cpp @@ -93,7 +93,7 @@ bool Qualifiers::isTargetAddressSpaceSupersetOf(LangAS A, LangAS B, // to implicitly cast into the default address space. (A == LangAS::Default && (B == LangAS::cuda_constant || B == LangAS::cuda_device || - B == LangAS::cuda_shared)) || + B == LangAS::cuda_shared || B == LangAS::amdgpu_barrier)) || // In HLSL, the this pointer for member functions points to the default // address space. This causes a problem if the structure is in // a different address space. We want to allow casting from these @@ -5492,6 +5492,31 @@ bool Type::isCUDADeviceBuiltinTextureType() const { return false; } +static bool isAMDGPUNamedBarrierTypeImpl(const Type *Ty, bool AllowWrappers) { + // This query does not care about qualifiers at all. + Ty = Ty->getUnqualifiedDesugaredType(); + + // Unwrap arrays. + while (isa<ArrayType>(Ty)) + Ty = Ty->getArrayElementTypeNoTypeQual()->getUnqualifiedDesugaredType(); + + if (const auto *BT = dyn_cast<BuiltinType>(Ty)) + return BT->getKind() == BuiltinType::AMDGPUNamedWorkgroupBarrier; + if (AllowWrappers) { + if (const auto *RT = dyn_cast<RecordType>(Ty)) + return RT->getDecl()->hasAttr<AMDGPUNamedBarrierWrapperAttr>(); + } + return false; +} + +bool Type::isAMDGPUNamedBarrierType() const { + return isAMDGPUNamedBarrierTypeImpl(this, /*AllowWrappers=*/false); +} + +bool Type::isAMDGPUNamedBarrierTypeOrWrapper() const { + return isAMDGPUNamedBarrierTypeImpl(this, /*AllowWrappers=*/true); +} + bool Type::hasSizedVLAType() const { if (!isVariablyModifiedType()) return false; diff --git a/clang/lib/AST/TypePrinter.cpp b/clang/lib/AST/TypePrinter.cpp index e8fbffb9f954d..aba901e764084 100644 --- a/clang/lib/AST/TypePrinter.cpp +++ b/clang/lib/AST/TypePrinter.cpp @@ -2748,6 +2748,8 @@ std::string Qualifiers::getAddrSpaceAsString(LangAS AS) { return "hlsl_push_constant"; case LangAS::wasm_funcref: return "__funcref"; + case LangAS::amdgpu_barrier: + return "amdgpu_barrier"; default: return std::to_string(toTargetAddressSpace(AS)); } diff --git a/clang/lib/Basic/TargetInfo.cpp b/clang/lib/Basic/TargetInfo.cpp index 46b5bdecb9c40..e0b745f8f23ac 100644 --- a/clang/lib/Basic/TargetInfo.cpp +++ b/clang/lib/Basic/TargetInfo.cpp @@ -21,6 +21,7 @@ #include "llvm/ADT/StringExtras.h" #include "llvm/Support/ErrorHandling.h" #include "llvm/TargetParser/TargetParser.h" +#include "llvm/TargetParser/Triple.h" #include <cstdlib> using namespace clang; @@ -55,6 +56,7 @@ static const LangASMap FakeAddrSpaceMap = { 18, // hlsl_output 19, // hlsl_push_constant 20, // wasm_funcref + 21, // amdgpu_barrier }; // TargetInfo Constructor. @@ -167,6 +169,7 @@ TargetInfo::TargetInfo(const llvm::Triple &T) : Triple(T) { HasBuiltinZOSVaList = false; HasAArch64ACLETypes = false; HasRISCVVTypes = false; + HasAMDGPUTypes = false; AllowAMDGPUUnsafeFPAtomics = false; HasUnalignedAccess = false; ARMCDECoprocMask = 0; diff --git a/clang/lib/Basic/Targets/AArch64.h b/clang/lib/Basic/Targets/AArch64.h index b6e707d8b4245..f02d17fa3cbac 100644 --- a/clang/lib/Basic/Targets/AArch64.h +++ b/clang/lib/Basic/Targets/AArch64.h @@ -54,6 +54,7 @@ static const unsigned ARM64AddrSpaceMap[] = { // Wasm address space values for this target are dummy values, // as it is only enabled for Wasm targets. 20, // wasm_funcref + 0, // amdgpu_barrier }; using AArch64FeatureSet = llvm::SmallDenseSet<StringRef, 32>; diff --git a/clang/lib/Basic/Targets/AMDGPU.cpp b/clang/lib/Basic/Targets/AMDGPU.cpp index 3fd9643373383..4665dac9cb787 100644 --- a/clang/lib/Basic/Targets/AMDGPU.cpp +++ b/clang/lib/Basic/Targets/AMDGPU.cpp @@ -56,6 +56,8 @@ const LangASMap AMDGPUTargetInfo::AMDGPUAddrSpaceMap = { llvm::AMDGPUAS::PRIVATE_ADDRESS, // hlsl_input llvm::AMDGPUAS::PRIVATE_ADDRESS, // hlsl_output llvm::AMDGPUAS::GLOBAL_ADDRESS, // hlsl_push_constant + llvm::AMDGPUAS::FLAT_ADDRESS, // wasm_funcref + llvm::AMDGPUAS::LOCAL_ADDRESS, // amdgpu_barrier }; } // namespace targets @@ -199,6 +201,7 @@ AMDGPUTargetInfo::AMDGPUTargetInfo(const llvm::Triple &Triple, AddrSpaceMap = &AMDGPUAddrSpaceMap; UseAddrSpaceMapMangling = true; + HasAMDGPUTypes = true; if (Triple.isAMDGCN()) { // __bf16 is always available as a load/store only type on AMDGCN. diff --git a/clang/lib/Basic/Targets/DirectX.h b/clang/lib/Basic/Targets/DirectX.h index 64e5533bbffeb..67725141eb47c 100644 --- a/clang/lib/Basic/Targets/DirectX.h +++ b/clang/lib/Basic/Targets/DirectX.h @@ -51,6 +51,7 @@ static const unsigned DirectXAddrSpaceMap[] = { // Wasm address space values for this target are dummy values, // as it is only enabled for Wasm targets. 20, // wasm_funcref + 0, // amdgpu_barrier }; class LLVM_LIBRARY_VISIBILITY DirectXTargetInfo : public TargetInfo { diff --git a/clang/lib/Basic/Targets/NVPTX.h b/clang/lib/Basic/Targets/NVPTX.h index 00be0fe54c1bd..803629e3406d4 100644 --- a/clang/lib/Basic/Targets/NVPTX.h +++ b/clang/lib/Basic/Targets/NVPTX.h @@ -55,6 +55,7 @@ static const unsigned NVPTXAddrSpaceMap[] = { // Wasm address space values for this target are dummy values, // as it is only enabled for Wasm targets. 20, // wasm_funcref + 0, // amdgpu_barrier }; /// The DWARF address class. Taken from diff --git a/clang/lib/Basic/Targets/SPIR.h b/clang/lib/Basic/Targets/SPIR.h index 42256920f353e..37baf7c99a90c 100644 --- a/clang/lib/Basic/Targets/SPIR.h +++ b/clang/lib/Basic/Targets/SPIR.h @@ -59,6 +59,7 @@ static const unsigned SPIRDefIsPrivMap[] = { // Wasm address space values for this target are dummy values, // as it is only enabled for Wasm targets. 20, // wasm_funcref + 0, // amdgpu_barrier }; // Used by both the SPIR and SPIR-V targets. @@ -97,6 +98,7 @@ static const unsigned SPIRDefIsGenMap[] = { // Wasm address space values for this target are dummy values, // as it is only enabled for Wasm targets. 20, // wasm_funcref + 0, // amdgpu_barrier }; // Base class for SPIR and SPIR-V target info. @@ -310,6 +312,7 @@ class LLVM_LIBRARY_VISIBILITY BaseSPIRVTargetInfo : public BaseSPIRTargetInfo { BaseSPIRVTargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts) : BaseSPIRTargetInfo(Triple, Opts) { assert(Triple.isSPIRV() && "Invalid architecture for SPIR-V."); + HasAMDGPUTypes = (Triple.getVendor() == llvm::Triple::AMD); } llvm::SmallVector<Builtin::InfosShard> getTargetBuiltins() const override; diff --git a/clang/lib/Basic/Targets/SystemZ.h b/clang/lib/Basic/Targets/SystemZ.h index c06b142200d75..c12cefc10e08a 100644 --- a/clang/lib/Basic/Targets/SystemZ.h +++ b/clang/lib/Basic/Targets/SystemZ.h @@ -48,7 +48,8 @@ static const unsigned ZOSAddressMap[] = { 0, // hlsl_input 0, // hlsl_output 0, // hlsl_push_constant - 0 // wasm_funcref + 0, // wasm_funcref + 0, // amdgpu_barrier }; class LLVM_LIBRARY_VISIBILITY SystemZTargetInfo : public TargetInfo { diff --git a/clang/lib/Basic/Targets/TCE.h b/clang/lib/Basic/Targets/TCE.h index 1360298de9794..403b8b1bb3304 100644 --- a/clang/lib/Basic/Targets/TCE.h +++ b/clang/lib/Basic/Targets/TCE.h @@ -60,6 +60,7 @@ static const unsigned TCEOpenCLAddrSpaceMap[] = { // Wasm address space values for this target are dummy values, // as it is only enabled for Wasm targets. 20, // wasm_funcref + 0, // amdgpu_barrier }; class LLVM_LIBRARY_VISIBILITY TCETargetInfo : public TargetInfo { diff --git a/clang/lib/Basic/Targets/WebAssembly.h b/clang/lib/Basic/Targets/WebAssembly.h index 0b0266a48469a..5389025817b02 100644 --- a/clang/lib/Basic/Targets/WebAssembly.h +++ b/clang/lib/Basic/Targets/WebAssembly.h @@ -49,6 +49,7 @@ static const unsigned WebAssemblyAddrSpaceMap[] = { 0, // hlsl_output 0, // hlsl_push_constant 20, // wasm_funcref + 0, // amdgpu_barrier }; class LLVM_LIBRARY_VISIBILITY WebAssemblyTargetInfo : public TargetInfo { diff --git a/clang/lib/Basic/Targets/X86.h b/clang/lib/Basic/Targets/X86.h index e305d9017d897..ab5cd9618bbaa 100644 --- a/clang/lib/Basic/Targets/X86.h +++ b/clang/lib/Basic/Targets/X86.h @@ -55,6 +55,7 @@ static const unsigned X86AddrSpaceMap[] = { // Wasm address space values for this target are dummy values, // as it is only enabled for Wasm targets. 20, // wasm_funcref + 0, // amdgpu_barrier }; // X86 target abstract base class; x86-32 and x86-64 are very close, so diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 78627047b19ad..6f32cf5cb5711 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -6220,6 +6220,9 @@ LangAS CodeGenModule::GetGlobalVarAddressSpace(const VarDecl *D) { if (LangOpts.CUDA && LangOpts.CUDAIsDevice) { if (D) { + if (D->getType()->isAMDGPUNamedBarrierTypeOrWrapper()) + return LangAS::amdgpu_barrier; + if (D->hasAttr<CUDAConstantAttr>()) return LangAS::cuda_constant; if (D->hasAttr<CUDASharedAttr>()) diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp index e689c75eb566c..07e83bc0c85ac 100644 --- a/clang/lib/Sema/Sema.cpp +++ b/clang/lib/Sema/Sema.cpp @@ -569,14 +569,9 @@ void Sema::Initialize() { #include "clang/Basic/WebAssemblyReferenceTypes.def" } - if (Context.getTargetInfo().getTriple().isAMDGPU() || - (Context.getTargetInfo().getTriple().isSPIRV() && - Context.getTargetInfo().getTriple().getVendor() == llvm::Triple::AMD) || + if (Context.getTargetInfo().hasAMDGPUTypes() || (Context.getAuxTargetInfo() && - (Context.getAuxTargetInfo()->getTriple().isAMDGPU() || - (Context.getAuxTargetInfo()->getTriple().isSPIRV() && - Context.getAuxTargetInfo()->getTriple().getVendor() == - llvm::Triple::AMD)))) { + (Context.getAuxTargetInfo()->hasAMDGPUTypes()))) { #define AMDGPU_TYPE(Name, Id, SingletonId, Width, Align) \ addImplicitTypedef(Name, Context.SingletonId); #include "clang/Basic/AMDGPUTypes.def" diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp index 48230fa262d5c..30718273718e7 100644 --- a/clang/lib/Sema/SemaAMDGPU.cpp +++ b/clang/lib/Sema/SemaAMDGPU.cpp @@ -1079,4 +1079,89 @@ bool DiagnoseUnguardedBuiltins::VisitCallExpr(CallExpr *CE) { void SemaAMDGPU::DiagnoseUnguardedBuiltinUsage(FunctionDecl *FD) { DiagnoseUnguardedBuiltins(SemaRef).IssueDiagnostics(FD->getBody()); } + +static FieldDecl *getNamedBarrierField(const RecordDecl *R) { + for (FieldDecl *FD : R->fields()) { + QualType FDTy = FD->getType(); + if (FDTy->isAMDGPUNamedBarrierTypeOrWrapper()) + return FD; + } + + return nullptr; +} + +void SemaAMDGPU::checkNamedBarrierWrapper(RecordDecl *R) { + ASTContext &Context = getASTContext(); + if (R->isInvalidDecl()) + return; + + if (!Context.getTargetInfo().hasAMDGPUTypes() && + (!Context.getAuxTargetInfo() || + !Context.getAuxTargetInfo()->hasAMDGPUTypes())) + return; + + bool IsWrapper = false; + std::function<void()> DiagWrapperNote; + + // First, check if this is a named barrier wrapper by virtue of the class + // declaring a named barrier field. This covers both C and C++. + if (FieldDecl *NamedBarrField = getNamedBarrierField(R)) { + // If this record contains a named barrier field, it must have only one + // field. + if (R->getNumFields() > 1) { + SemaRef.Diag(NamedBarrField->getLocation(), + diag::err_amdgcn_invalid_field_not_a_wrapper) + << NamedBarrField->getType(); + SemaRef.Diag( + R->getLocation(), + diag::note_amdgcn_not_a_named_barrier_wrapper_too_many_fields) + << R->getDeclName(); + return; + } + + IsWrapper = true; + DiagWrapperNote = [this, R, NamedBarrField]() { + SemaRef.Diag(NamedBarrField->getLocation(), + diag::note_amdgcn_named_barrier_reason_field) + << R->getDeclName() << NamedBarrField->getDeclName(); + }; + } + + // Then, for C++ classes, check if this is a named barrier wrapper by virtue + // of inheriting one. + const auto *CxxR = dyn_cast<CXXRecordDecl>(R); + if (CxxR && !IsWrapper) { + for (CXXBaseSpecifier BS : CxxR->bases()) { + const RecordDecl *Base = BS.getType()->getAsRecordDecl(); + if (!Base || !Base->hasAttr<AMDGPUNamedBarrierWrapperAttr>()) + continue; + + IsWrapper = true; + DiagWrapperNote = [this, BS, R]() { + // Print using the CXXBaseSpecifier type as it includes the template + // parameters. + SemaRef.Diag(BS.getBeginLoc(), + diag::note_amdgcn_named_barrier_reason_inherited) + << R->getDeclName() << BS.getType(); + }; + } + } + + if (!IsWrapper) + return; + + // Set the attribute even if the wrapper may be found to be invalid later. + R->addAttr( + AMDGPUNamedBarrierWrapperAttr::CreateImplicit(Context, SourceRange())); + + // This is a wrapper CXXRecordDecl, it must have a C++11 standard layout. + if (CxxR && !CxxR->isCXX11StandardLayout()) { + SemaRef.Diag(R->getLocation(), + diag::err_amdgcn_named_barrier_wrapper_non_standard_layout) + << R->getDeclName(); + assert(DiagWrapperNote && + "IsWrapper is set but no context diagnostic provided"); + DiagWrapperNote(); + } +} } // namespace clang diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 7e1b23c971a9c..2d1ecb163bb18 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -20466,6 +20466,10 @@ void Sema::ActOnFields(Scope *S, SourceLocation RecLoc, Decl *EnclosingDecl, CDecl->setIvarRBraceLoc(RBrac); } } + + if (Record) + AMDGPU().checkNamedBarrierWrapper(Record); + if (Record && !isa<ClassTemplateSpecializationDecl>(Record)) ProcessAPINotes(Record); } diff --git a/clang/test/CodeGenHIP/amdgpu-barrier-type.hip b/clang/test/CodeGenHIP/amdgpu-barrier-type.hip index 947ceb56d279e..df9e3631c0d1f 100644 --- a/clang/test/CodeGenHIP/amdgpu-barrier-type.hip +++ b/clang/test/CodeGenHIP/amdgpu-barrier-type.hip @@ -1,18 +1,28 @@ -// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature - // REQUIRES: amdgpu-registered-target - // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu verde -emit-llvm -o - %s | FileCheck %s +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --check-globals --global-value-regex "bar.*" +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -target-cpu gfx1250 -emit-llvm -o - %s | FileCheck %s #define __shared__ __attribute__((shared)) __shared__ __amdgpu_named_workgroup_barrier_t bar; -__shared__ __amdgpu_named_workgroup_barrier_t arr[2]; -__shared__ struct { +__shared__ __amdgpu_named_workgroup_barrier_t bar_arr[2]; + +//. +// CHECK: @bar = addrspace(3) global target("amdgcn.named.barrier", 0) undef, align 4 +// CHECK: @bar_arr = addrspace(3) global [2 x target("amdgcn.named.barrier", 0)] undef, align 4 +// CHECK: @bar_wrapper_str = addrspace(3) global %struct.WrapperStruct undef, align 4 +// CHECK: @bar_wrapperwrapper_str = addrspace(3) global %struct.WrapperWrapperStruct undef, align 4 +//. +__shared__ struct WrapperStruct { __amdgpu_named_workgroup_barrier_t x; - __amdgpu_named_workgroup_barrier_t y; -} str; +} bar_wrapper_str; + +__shared__ struct WrapperWrapperStruct { + WrapperStruct x; +} bar_wrapperwrapper_str; -__amdgpu_named_workgroup_barrier_t *getBar(); -void useBar(__amdgpu_named_workgroup_barrier_t *); +__attribute__((device)) __amdgpu_named_workgroup_barrier_t *getBar(); +__attribute__((device)) void useBar(__amdgpu_named_workgroup_barrier_t *); // CHECK-LABEL: define {{[^@]+}}@_Z7testSemPu34__amdgpu_named_workgroup_barrier_t // CHECK-SAME: (ptr noundef [[P:%.*]]) #[[ATTR0:[0-9]+]] { @@ -22,19 +32,21 @@ void useBar(__amdgpu_named_workgroup_barrier_t *); // CHECK-NEXT: store ptr [[P]], ptr [[P_ADDR_ASCAST]], align 8 // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[P_ADDR_ASCAST]], align 8 // CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef [[TMP0]]) #[[ATTR2:[0-9]+]] -// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef addrspacecast (ptr addrspace(1) @bar to ptr)) #[[ATTR2]] -// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef getelementptr inbounds nuw (i8, ptr addrspacecast (ptr addrspace(1) @arr to ptr), i64 16)) #[[ATTR2]] -// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef getelementptr inbounds nuw (i8, ptr addrspacecast (ptr addrspace(1) @str to ptr), i64 16)) #[[ATTR2]] +// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef addrspacecast (ptr addrspace(3) @bar to ptr)) #[[ATTR2]] +// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef addrspacecast (ptr addrspace(3) @bar_wrapper_str to ptr)) #[[ATTR2]] +// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef addrspacecast (ptr addrspace(3) @bar_wrapperwrapper_str to ptr)) #[[ATTR2]] +// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef getelementptr inbounds nuw (i8, ptr addrspacecast (ptr addrspace(3) @bar_arr to ptr), i64 16)) #[[ATTR2]] // CHECK-NEXT: [[CALL:%.*]] = call noundef ptr @_Z6getBarv() #[[ATTR2]] // CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef [[CALL]]) #[[ATTR2]] // CHECK-NEXT: [[CALL1:%.*]] = call noundef ptr @_Z6getBarv() #[[ATTR2]] // CHECK-NEXT: ret ptr [[CALL1]] // -__amdgpu_named_workgroup_barrier_t *testSem(__amdgpu_named_workgroup_barrier_t *p) { +__attribute__((device)) __amdgpu_named_workgroup_barrier_t *testSem(__amdgpu_named_workgroup_barrier_t *p) { useBar(p); useBar(&bar); - useBar(&arr[1]); - useBar(&str.y); + useBar(&bar_wrapper_str.x); + useBar(&bar_wrapperwrapper_str.x.x); + useBar(&bar_arr[1]); useBar(getBar()); return getBar(); } diff --git a/clang/test/SemaCXX/amdgpu-barrier.cpp b/clang/test/SemaCXX/amdgpu-barrier.cpp index a171433727dda..80816fba6105e 100644 --- a/clang/test/SemaCXX/amdgpu-barrier.cpp +++ b/clang/test/SemaCXX/amdgpu-barrier.cpp @@ -13,5 +13,76 @@ void foo() { void *vp = (void *)k; // expected-error {{cannot cast from type '__amdgpu_named_workgroup_barrier_t' to pointer type 'void *'}} } +using SugaredArray = __amdgpu_named_workgroup_barrier_t[2]; + +struct TestSimple { + __amdgpu_named_workgroup_barrier_t x; +}; + +struct TestArray{ + __amdgpu_named_workgroup_barrier_t y[2]; +}; + +struct TestSugared { + SugaredArray z[2]; +}; + +struct TestSimpleWithStaticField { + __amdgpu_named_workgroup_barrier_t x; + static unsigned Harmless; +}; + +// Wrappers cannot have >1 field. +struct WrapperHasTooManyFields { // expected-note {{'WrapperHasTooManyFields' is not a named barrier wrapper because it has more than one field}} + __amdgpu_named_workgroup_barrier_t x; // expected-error {{fields of type '__amdgpu_named_workgroup_barrier_t' are only allowed in named barrier wrappers}} + int other; +}; + +// Wrappers must have standard layout +struct WrapperBase { +__amdgpu_named_workgroup_barrier_t x; +}; + +struct WrapperWithBase : public WrapperBase { +}; + +// expected-error@+2 {{named barrier wrapper 'WrapperWithBaseNoStandardLayout' must have a C++11 standard layou}} +// expected-note@+1 {{'WrapperWithBaseNoStandardLayout' is a named barrier wrapper because it inherits from named barrier wrapper 'WrapperBase'}} +struct WrapperWithBaseNoStandardLayout : public WrapperBase { + unsigned K = 0; +}; + +// Wrappers of Wrappers have the same restrictions. +struct WrapperOfWrapperWithTooManyFields { // expected-note {{'WrapperOfWrapperWithTooManyFields' is not a named barrier wrapper because it has more than one field}} + TestSimple x; // expected-error {{fields of type 'TestSimple' are only allowed in named barrier wrappers}} + int other; +}; + +struct WrapperOfWrapper { + WrapperWithBase y; +}; + +// Check templated cases with a bit of complexity thrown in. +template<typename Derived> +class CRTPWrapperBase { + CRTPWrapperBase() { + static_cast<Derived*>(this)->sayHello(); + } + + WrapperOfWrapper wow; +}; + +class TemplatedWrapperImpl : public CRTPWrapperBase<TemplatedWrapperImpl> { + void sayHello() {} +}; + +// expected-error@+2 {{named barrier wrapper 'TemplatedWrapperImplNoStandardLayout' must have a C++11 standard layou}} +// expected-note@+1 {{'TemplatedWrapperImplNoStandardLayout' is a named barrier wrapper because it inherits from named barrier wrapper 'CRTPWrapperBase<TemplatedWrapperImpl>'}} +class TemplatedWrapperImplNoStandardLayout : public CRTPWrapperBase<TemplatedWrapperImpl> { + void sayHello() {} + + int k = 0; +}; + static_assert(sizeof(__amdgpu_named_workgroup_barrier_t) == 16, "wrong size"); static_assert(alignof(__amdgpu_named_workgroup_barrier_t) == 4, "wrong alignment"); diff --git a/clang/test/SemaHIP/amdgpu-barrier.hip b/clang/test/SemaHIP/amdgpu-barrier.hip index ccd99b1e2c1f2..e8cd52d77d4a5 100644 --- a/clang/test/SemaHIP/amdgpu-barrier.hip +++ b/clang/test/SemaHIP/amdgpu-barrier.hip @@ -16,5 +16,77 @@ __device__ void foo() { void *vp = (void *)k; // expected-error {{cannot cast from type '__amdgpu_named_workgroup_barrier_t' to pointer type 'void *'}} } +using SugaredArray = __amdgpu_named_workgroup_barrier_t[2]; + +struct TestSimple { + __amdgpu_named_workgroup_barrier_t x; +}; + +struct TestArray{ + __amdgpu_named_workgroup_barrier_t y[2]; +}; + +struct TestSugared { + SugaredArray z[2]; +}; + +struct TestSimpleWithStaticField { + __amdgpu_named_workgroup_barrier_t x; + static unsigned Harmless; +}; + +// Wrappers cannot have >1 field. +struct WrapperHasTooManyFields { // expected-note {{'WrapperHasTooManyFields' is not a named barrier wrapper because it has more than one field}} + __amdgpu_named_workgroup_barrier_t x; // expected-error {{fields of type '__amdgpu_named_workgroup_barrier_t' are only allowed in named barrier wrappers}} + int other; +}; + +// Wrappers must have standard layout +struct WrapperBase { +__amdgpu_named_workgroup_barrier_t x; +}; + +struct WrapperWithBase : public WrapperBase { +}; + +// expected-error@+2 {{named barrier wrapper 'WrapperWithBaseNoStandardLayout' must have a C++11 standard layou}} +// expected-note@+1 {{'WrapperWithBaseNoStandardLayout' is a named barrier wrapper because it inherits from named barrier wrapper 'WrapperBase'}} +struct WrapperWithBaseNoStandardLayout : public WrapperBase { + unsigned K = 0; +}; + +// Wrappers of Wrappers have the same restrictions. +struct WrapperOfWrapperWithTooManyFields { // expected-note {{'WrapperOfWrapperWithTooManyFields' is not a named barrier wrapper because it has more than one field}} + TestSimple x; // expected-error {{fields of type 'TestSimple' are only allowed in named barrier wrappers}} + int other; +}; + +struct WrapperOfWrapper { + WrapperWithBase y; +}; + +// Check templated cases with a bit of complexity thrown in. +template<typename Derived> +class CRTPWrapperBase { + CRTPWrapperBase() { + static_cast<Derived*>(this)->sayHello(); + } + + WrapperOfWrapper wow; +}; + +class TemplatedWrapperImpl : public CRTPWrapperBase<TemplatedWrapperImpl> { + void sayHello() {} +}; + +// expected-error@+2 {{named barrier wrapper 'TemplatedWrapperImplNoStandardLayout' must have a C++11 standard layou}} +// expected-note@+1 {{'TemplatedWrapperImplNoStandardLayout' is a named barrier wrapper because it inherits from named barrier wrapper 'CRTPWrapperBase<TemplatedWrapperImpl>'}} +class TemplatedWrapperImplNoStandardLayout : public CRTPWrapperBase<TemplatedWrapperImpl> { + void sayHello() {} + + int k = 0; +}; + + static_assert(sizeof(__amdgpu_named_workgroup_barrier_t) == 16, "wrong size"); static_assert(alignof(__amdgpu_named_workgroup_barrier_t) == 4, "wrong alignment"); diff --git a/clang/test/SemaOpenCL/amdgpu-barrier.cl b/clang/test/SemaOpenCL/amdgpu-barrier.cl index 150c311c7c593..0898b1642f96f 100644 --- a/clang/test/SemaOpenCL/amdgpu-barrier.cl +++ b/clang/test/SemaOpenCL/amdgpu-barrier.cl @@ -3,6 +3,30 @@ // RUN: %clang_cc1 -verify -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -Wno-unused-value %s void foo() { + typedef __amdgpu_named_workgroup_barrier_t SugaredArray[2]; + + struct TestSimple { + __amdgpu_named_workgroup_barrier_t x; + }; + + struct TestArray { + __amdgpu_named_workgroup_barrier_t y[2]; + }; + + struct TestSugared { + SugaredArray z[2]; + }; + + struct GoodWrapper { + __amdgpu_named_workgroup_barrier_t x; + }; + + // Wrappers cannot have >1 field. + struct WrapperHasTooManyFields { // expected-note {{'WrapperHasTooManyFields' is not a named barrier wrapper because it has more than one field}} + __amdgpu_named_workgroup_barrier_t x; // expected-error {{fields of type '__amdgpu_named_workgroup_barrier_t' are only allowed in named barrier wrappers}} + int other; + }; + int n = 100; __amdgpu_named_workgroup_barrier_t v = 0; // expected-error {{initializing '__private __amdgpu_named_workgroup_barrier_t' with an expression of incompatible type 'int'}} int c = v; // expected-error {{initializing '__private int' with an expression of incompatible type '__private __amdgpu_named_workgroup_barrier_t'}} diff --git a/clang/test/SemaTemplate/address_space-dependent.cpp b/clang/test/SemaTemplate/address_space-dependent.cpp index 3fdccb2c71a76..d6f25923b69b5 100644 --- a/clang/test/SemaTemplate/address_space-dependent.cpp +++ b/clang/test/SemaTemplate/address_space-dependent.cpp @@ -43,7 +43,7 @@ void neg() { template <long int I> void tooBig() { - __attribute__((address_space(I))) int *bounds; // expected-error {{address space is larger than the maximum supported (8388580)}} + __attribute__((address_space(I))) int *bounds; // expected-error {{address space is larger than the maximum supported (8388579)}} } template <long int I> @@ -101,7 +101,7 @@ int main() { car<1, 2, 3>(); // expected-note {{in instantiation of function template specialization 'car<1, 2, 3>' requested here}} HasASTemplateFields<1> HASTF; neg<-1>(); // expected-note {{in instantiation of function template specialization 'neg<-1>' requested here}} - correct<0x7FFFE4>(); + correct<0x7FFFE3>(); tooBig<8388650>(); // expected-note {{in instantiation of function template specialization 'tooBig<8388650L>' requested here}} __attribute__((address_space(1))) char *x; _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
