https://github.com/Pierre-vh updated https://github.com/llvm/llvm-project/pull/207687
>From f605bb016078a0fdc061562e2a3338b0204806e6 Mon Sep 17 00:00:00 2001 From: pvanhout <[email protected]> Date: Wed, 15 Jul 2026 13:04:14 +0200 Subject: [PATCH 1/5] [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/AMDGPU.cpp | 2 + clang/lib/Basic/Targets/SPIR.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 +- 21 files changed, 365 insertions(+), 33 deletions(-) diff --git a/clang/include/clang/AST/TypeBase.h b/clang/include/clang/AST/TypeBase.h index 00b6f1b20ad8e..a0a470803d68d 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 2dfaa1c45ac55..1833dd1d4a4a8 100644 --- a/clang/include/clang/Basic/AddressSpaces.h +++ b/clang/include/clang/Basic/AddressSpaces.h @@ -71,6 +71,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 39c672322d515..ff8c8160ad6c1 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">]; @@ -5333,7 +5340,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 89e2f956971b3..cf549171c2049 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -14274,6 +14274,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 968a0c1b129ef..6311b6b567a5e 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 40daf646a4711..474c67612fd1c 100644 --- a/clang/lib/AST/TypePrinter.cpp +++ b/clang/lib/AST/TypePrinter.cpp @@ -2751,6 +2751,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 103104ce4874b..1bb0026d8422e 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 constexpr LangASMap FakeAddrSpaceMap = { {LangAS::hlsl_output, 18}, {LangAS::hlsl_push_constant, 19}, {LangAS::wasm_funcref, 20}, + {LangAS::amdgpu_barrier, 21}, }; // 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/AMDGPU.cpp b/clang/lib/Basic/Targets/AMDGPU.cpp index c789fd8f94afb..285153695da27 100644 --- a/clang/lib/Basic/Targets/AMDGPU.cpp +++ b/clang/lib/Basic/Targets/AMDGPU.cpp @@ -56,6 +56,7 @@ const LangASMap AMDGPUTargetInfo::AMDGPUAddrSpaceMap = { {LangAS::hlsl_input, llvm::AMDGPUAS::PRIVATE_ADDRESS}, {LangAS::hlsl_output, llvm::AMDGPUAS::PRIVATE_ADDRESS}, {LangAS::hlsl_push_constant, llvm::AMDGPUAS::GLOBAL_ADDRESS}, + {LangAS::amdgpu_barrier, llvm::AMDGPUAS::LOCAL_ADDRESS}, }; } // namespace targets @@ -199,6 +200,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/SPIR.h b/clang/lib/Basic/Targets/SPIR.h index a16c8ec79d2f2..a5d58dcf2c835 100644 --- a/clang/lib/Basic/Targets/SPIR.h +++ b/clang/lib/Basic/Targets/SPIR.h @@ -278,6 +278,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/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 817b2b5e6a69d..4d534ec713c36 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -6222,6 +6222,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 9f962912148ab..a63897add6403 100644 --- a/clang/lib/Sema/Sema.cpp +++ b/clang/lib/Sema/Sema.cpp @@ -568,14 +568,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 7de5542e72559..f3614045fbebb 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -20488,6 +20488,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; >From b624029c6e0a61dd2db8e0a6f5de12bb58301149 Mon Sep 17 00:00:00 2001 From: pvanhout <[email protected]> Date: Thu, 16 Jul 2026 14:32:18 +0200 Subject: [PATCH 2/5] Comment --- clang/include/clang/Basic/DiagnosticSemaKinds.td | 2 +- clang/test/SemaCXX/amdgpu-barrier.cpp | 4 ++-- clang/test/SemaHIP/amdgpu-barrier.hip | 4 ++-- clang/test/SemaOpenCL/amdgpu-barrier.cl | 2 +- 4 files changed, 6 insertions(+), 6 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index cf549171c2049..d7531e4195848 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -14276,7 +14276,7 @@ def note_acc_reduction_combiner_forming // AMDGCN type diagnostics def err_amdgcn_invalid_field_not_a_wrapper : Error< - "fields of type %0 are only allowed in named barrier wrappers">; + "field with barrier type %0 seen in a structure that is not a named barrier wrapper">; 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">; diff --git a/clang/test/SemaCXX/amdgpu-barrier.cpp b/clang/test/SemaCXX/amdgpu-barrier.cpp index 80816fba6105e..b7e376624a268 100644 --- a/clang/test/SemaCXX/amdgpu-barrier.cpp +++ b/clang/test/SemaCXX/amdgpu-barrier.cpp @@ -34,7 +34,7 @@ struct TestSimpleWithStaticField { // 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}} + __amdgpu_named_workgroup_barrier_t x; // expected-error {{field with barrier type '__amdgpu_named_workgroup_barrier_t' seen in a structure that is not a named barrier wrapper}} int other; }; @@ -54,7 +54,7 @@ struct WrapperWithBaseNoStandardLayout : public WrapperBase { // 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}} + TestSimple x; // expected-error {{field with barrier type 'TestSimple' seen in a structure that is not a named barrier wrapper}} int other; }; diff --git a/clang/test/SemaHIP/amdgpu-barrier.hip b/clang/test/SemaHIP/amdgpu-barrier.hip index e8cd52d77d4a5..e8abfcaf1df84 100644 --- a/clang/test/SemaHIP/amdgpu-barrier.hip +++ b/clang/test/SemaHIP/amdgpu-barrier.hip @@ -37,7 +37,7 @@ struct TestSimpleWithStaticField { // 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}} + __amdgpu_named_workgroup_barrier_t x; // expected-error {{field with barrier type '__amdgpu_named_workgroup_barrier_t' seen in a structure that is not a named barrier wrapper}} int other; }; @@ -57,7 +57,7 @@ struct WrapperWithBaseNoStandardLayout : public WrapperBase { // 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}} + TestSimple x; // expected-error {{field with barrier type 'TestSimple' seen in a structure that is not a named barrier wrapper}} int other; }; diff --git a/clang/test/SemaOpenCL/amdgpu-barrier.cl b/clang/test/SemaOpenCL/amdgpu-barrier.cl index 0898b1642f96f..15e5d447c91f6 100644 --- a/clang/test/SemaOpenCL/amdgpu-barrier.cl +++ b/clang/test/SemaOpenCL/amdgpu-barrier.cl @@ -23,7 +23,7 @@ void foo() { // 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}} + __amdgpu_named_workgroup_barrier_t x; // expected-error {{field with barrier type '__amdgpu_named_workgroup_barrier_t' seen in a structure that is not a named barrier wrapper}} int other; }; >From 5142852da7f1f17687509c34514b4a32765ccf71 Mon Sep 17 00:00:00 2001 From: pvanhout <[email protected]> Date: Wed, 22 Jul 2026 10:28:10 +0200 Subject: [PATCH 3/5] Add docs --- clang/docs/AMDGPUSupport.md | 27 +++++++++++++++++++++++++++ 1 file changed, 27 insertions(+) diff --git a/clang/docs/AMDGPUSupport.md b/clang/docs/AMDGPUSupport.md index 8731af371b23f..84d0bb16e77c2 100644 --- a/clang/docs/AMDGPUSupport.md +++ b/clang/docs/AMDGPUSupport.md @@ -30,3 +30,30 @@ Please note that the specific architecture and feature names will vary depending Clang exposes AMDGPU hardware intrinsics as target-specific builtins with the `__builtin_amdgcn_` prefix. These are documented in {doc}`AMDGPUBuiltinReference`. + +## Target-Specific Types + +### Named Workgroup barrier Type + +The `__amdgpu_named_workgroup_barrier_t` type is used to represent the GFX12.5 named barriers. +Example usage: + +```c +__amdgpu_named_workgroup_barrier_t x; +__amdgpu_named_workgroup_barrier_t arr[2]; // Arrays are also fine + +void foo(int a) +{ + __builtin_amdgcn_s_barrier_init(&x, a); +} +``` + +When a class has a field of this type, the entire class is considered as a +"named barrier wrapper". Named barrier wrappers, and any derived types, are subject to the +following limitations: + +* They must be a standard-layout type (see `std:;is_standard_layout`). +* They have at most one field. + +When a class has a field that is a named barrier wrapper, the same restrictions also +apply and that class is also considered a named barrier wrapper. >From 620f7a6336a4691d475a77786d44566089a301ed Mon Sep 17 00:00:00 2001 From: pvanhout <[email protected]> Date: Wed, 22 Jul 2026 16:53:47 +0200 Subject: [PATCH 4/5] docs --- clang/docs/AMDGPUSupport.md | 24 ++++++++++++++++-------- 1 file changed, 16 insertions(+), 8 deletions(-) diff --git a/clang/docs/AMDGPUSupport.md b/clang/docs/AMDGPUSupport.md index 84d0bb16e77c2..f75e0b001f015 100644 --- a/clang/docs/AMDGPUSupport.md +++ b/clang/docs/AMDGPUSupport.md @@ -33,7 +33,7 @@ Clang exposes AMDGPU hardware intrinsics as target-specific builtins with the ## Target-Specific Types -### Named Workgroup barrier Type +### Named Workgroup Barrier Type The `__amdgpu_named_workgroup_barrier_t` type is used to represent the GFX12.5 named barriers. Example usage: @@ -48,12 +48,20 @@ void foo(int a) } ``` -When a class has a field of this type, the entire class is considered as a -"named barrier wrapper". Named barrier wrappers, and any derived types, are subject to the -following limitations: +A "named barrier wrapper" is a class that contains exactly one field, which is either +a single value or an array of values of one of the following types: -* They must be a standard-layout type (see `std:;is_standard_layout`). -* They have at most one field. +* `__amdgpu_named_workgroup_barrier_t`. +* Another "named barrier wrapper". -When a class has a field that is a named barrier wrapper, the same restrictions also -apply and that class is also considered a named barrier wrapper. +Named barrier wrappers let users add helper methods around named barrier objects. + +In C++, a class that inherits from a named barrier wrapper is also considered a +named barrier wrapper. Named barrier wrappers must be standard-layout types +(see `std::is_standard_layout`). +This means that named barrier wrappers: + +* May not have a virtual table: they cannot declare or inherit any virtual + functions, or inherit from a virtual base class. +* May not have any extra fields, either declared by the class or inherited + from a base class. >From f0d6b5f0d253041e26c9bbbcaae4f1fb291cf6c7 Mon Sep 17 00:00:00 2001 From: pvanhout <[email protected]> Date: Fri, 24 Jul 2026 10:44:56 +0200 Subject: [PATCH 5/5] comments --- clang/docs/AMDGPUSupport.md | 24 ++++++++++++++++-------- 1 file changed, 16 insertions(+), 8 deletions(-) diff --git a/clang/docs/AMDGPUSupport.md b/clang/docs/AMDGPUSupport.md index f75e0b001f015..3306242d9d7e6 100644 --- a/clang/docs/AMDGPUSupport.md +++ b/clang/docs/AMDGPUSupport.md @@ -36,30 +36,38 @@ Clang exposes AMDGPU hardware intrinsics as target-specific builtins with the ### Named Workgroup Barrier Type The `__amdgpu_named_workgroup_barrier_t` type is used to represent the GFX12.5 named barriers. + +This type is subject to certain restrictions when used to declare non-static fields of classes. +See the {ref}`named barrier wrappers<namedbarrierwrappers>` section below for more information. + Example usage: ```c __amdgpu_named_workgroup_barrier_t x; __amdgpu_named_workgroup_barrier_t arr[2]; // Arrays are also fine +struct SimpleWrapper { + __amdgpu_named_workgroup_barrier_t foo; +}; + void foo(int a) { __builtin_amdgcn_s_barrier_init(&x, a); } ``` -A "named barrier wrapper" is a class that contains exactly one field, which is either -a single value or an array of values of one of the following types: +(namedbarrierwrappers)= +### Named Barrier Wrappers -* `__amdgpu_named_workgroup_barrier_t`. -* Another "named barrier wrapper". +A "named barrier wrapper" is a class that contains exactly one non-static field +of one of the following types: -Named barrier wrappers let users add helper methods around named barrier objects. +* `__amdgpu_named_workgroup_barrier_t`, or an array thereof. +* Another named barrier wrapper, or an array thereof. In C++, a class that inherits from a named barrier wrapper is also considered a -named barrier wrapper. Named barrier wrappers must be standard-layout types -(see `std::is_standard_layout`). -This means that named barrier wrappers: +named barrier wrapper, and all named barrier wrappers must be standard-layout +types (see `std::is_standard_layout`). This means that named barrier wrappers: * May not have a virtual table: they cannot declare or inherit any virtual functions, or inherit from a virtual base class. _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
