https://github.com/Pierre-vh updated https://github.com/llvm/llvm-project/pull/207687
>From 19b08ca9deda56689f1db6e27db0eacb4c9fc819 Mon Sep 17 00:00:00 2001 From: pvanhout <[email protected]> Date: Mon, 6 Jul 2026 11:33:34 +0200 Subject: [PATCH 01/10] [clang][AMDGPU] Clean-up handling of named barrier type - Do not allow the type in struct fields. This is more like a handle/resource than a real type. It does not follow the traditional C++ object model, and using it in a struct field can do some weird things if you instantiate too many of them. - Use a `hip_barrier` LangAS for this type that currently maps to the local AS. This allows easy switching to the barrier AS in a future patch. Alternative to #195612, see also #195613 --- clang/include/clang/AST/TypeBase.h | 3 ++ clang/include/clang/Basic/AddressSpaces.h | 3 ++ .../clang/Basic/DiagnosticSemaKinds.td | 2 +- clang/lib/AST/Type.cpp | 14 ++++++- clang/lib/AST/TypePrinter.cpp | 2 + clang/lib/Basic/TargetInfo.cpp | 1 + clang/lib/Basic/Targets/AArch64.h | 1 + clang/lib/Basic/Targets/AMDGPU.cpp | 2 + clang/lib/Basic/Targets/DirectX.h | 1 + clang/lib/Basic/Targets/NVPTX.h | 1 + clang/lib/Basic/Targets/SPIR.h | 2 + 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 | 6 +++ clang/lib/Sema/SemaDecl.cpp | 10 ++++- clang/test/CodeGenHIP/amdgpu-barrier-type.hip | 37 ++++++++++++------- clang/test/SemaCXX/amdgpu-barrier.cpp | 4 ++ clang/test/SemaHIP/amdgpu-barrier.hip | 5 +++ clang/test/SemaOpenCL/amdgpu-barrier.cl | 5 +++ .../SemaTemplate/address_space-dependent.cpp | 4 +- 22 files changed, 89 insertions(+), 20 deletions(-) diff --git a/clang/include/clang/AST/TypeBase.h b/clang/include/clang/AST/TypeBase.h index c9658775f0470..2a3aa9ebe4887 100644 --- a/clang/include/clang/AST/TypeBase.h +++ b/clang/include/clang/AST/TypeBase.h @@ -2813,6 +2813,9 @@ 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. + bool isAMDGPUNamedBarrierType() 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..d16fc8069ae43 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 + hip_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/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 3ff7e30d9f1f7..e378fe8750718 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -11885,7 +11885,7 @@ def note_within_param_of_type : Note< "within parameter %0 of type %1 declared here">; def note_illegal_field_declared_here : Note< "field of illegal %select{type|pointer type}0 %1 declared here">; -def err_opencl_type_struct_or_union_field : Error< +def err_invalid_type_for_struct_or_union_field : Error< "the %0 type cannot be used to declare a structure or union field">; def err_event_t_addr_space_qual : Error< "the event_t type can only be used with __private address space qualifier">; diff --git a/clang/lib/AST/Type.cpp b/clang/lib/AST/Type.cpp index 42d148715bc40..cb992eac05518 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::hip_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,18 @@ bool Type::isCUDADeviceBuiltinTextureType() const { return false; } +bool Type::isAMDGPUNamedBarrierType() const { + const Type *Ty = getUnqualifiedDesugaredType(); + + // unwrap arrays + while (isa<ArrayType>(Ty)) + Ty = Ty->getArrayElementTypeNoTypeQual(); + + if (const auto *BT = dyn_cast<BuiltinType>(Ty->getUnqualifiedDesugaredType())) + return BT->getKind() == BuiltinType::AMDGPUNamedWorkgroupBarrier; + return false; +} + bool Type::hasSizedVLAType() const { if (!isVariablyModifiedType()) return false; diff --git a/clang/lib/AST/TypePrinter.cpp b/clang/lib/AST/TypePrinter.cpp index e8fbffb9f954d..e467397d9df52 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::hip_barrier: + return "hip_barrier"; default: return std::to_string(toTargetAddressSpace(AS)); } diff --git a/clang/lib/Basic/TargetInfo.cpp b/clang/lib/Basic/TargetInfo.cpp index 46b5bdecb9c40..95cbe66ec54d5 100644 --- a/clang/lib/Basic/TargetInfo.cpp +++ b/clang/lib/Basic/TargetInfo.cpp @@ -55,6 +55,7 @@ static const LangASMap FakeAddrSpaceMap = { 18, // hlsl_output 19, // hlsl_push_constant 20, // wasm_funcref + 21, // hip_barrier }; // TargetInfo Constructor. diff --git a/clang/lib/Basic/Targets/AArch64.h b/clang/lib/Basic/Targets/AArch64.h index b6e707d8b4245..ccf5333370925 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, // hip_barrier }; using AArch64FeatureSet = llvm::SmallDenseSet<StringRef, 32>; diff --git a/clang/lib/Basic/Targets/AMDGPU.cpp b/clang/lib/Basic/Targets/AMDGPU.cpp index 50f9c1aa1aa02..1619d2bbfb8b7 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, // hip_barrier }; } // namespace targets diff --git a/clang/lib/Basic/Targets/DirectX.h b/clang/lib/Basic/Targets/DirectX.h index 64e5533bbffeb..1b6785d92d3e3 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, // hip_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..534d88e5e333a 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, // hip_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..ae2b598b46d67 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, // hip_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, // hip_barrier }; // Base class for SPIR and SPIR-V target info. diff --git a/clang/lib/Basic/Targets/SystemZ.h b/clang/lib/Basic/Targets/SystemZ.h index c06b142200d75..6d1420afc58ca 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, // hip_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..95a17e0a9356f 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, // hip_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..e991e1bd19d6f 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, // hip_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..6a52c80662d3b 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, // hip_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 012a1115ca7d3..ed76ea5c53b8a 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -6227,6 +6227,12 @@ LangAS CodeGenModule::GetGlobalVarAddressSpace(const VarDecl *D) { if (LangOpts.CUDA && LangOpts.CUDAIsDevice) { if (D) { + // TOOD: Forbid type on struct fields + // llvm::dbgs() << D->getNameAsString() << ": " << + // D->getType()->isAMDGPUNamedBarrierType() << "\n"; + if (D->getType()->isAMDGPUNamedBarrierType()) + return LangAS::hip_barrier; + if (D->hasAttr<CUDAConstantAttr>()) return LangAS::cuda_constant; if (D->hasAttr<CUDASharedAttr>()) diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 1d5a98f96d4ea..2dd524611423a 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -19476,7 +19476,7 @@ FieldDecl *Sema::CheckFieldDecl(DeclarationName Name, QualType T, // used as structure or union field: image, sampler, event or block types. if (T->isEventT() || T->isImageType() || T->isSamplerT() || T->isBlockPointerType()) { - Diag(Loc, diag::err_opencl_type_struct_or_union_field) << T; + Diag(Loc, diag::err_invalid_type_for_struct_or_union_field) << T; Record->setInvalidDecl(); InvalidDecl = true; } @@ -19489,6 +19489,14 @@ FieldDecl *Sema::CheckFieldDecl(DeclarationName Name, QualType T, } } + // AMDGPU does not allows the following types to be used for structure or + // union fields: named barriers + if (T->isAMDGPUNamedBarrierType()) { + Diag(Loc, diag::err_invalid_type_for_struct_or_union_field) << T; + Record->setInvalidDecl(); + InvalidDecl = true; + } + // Anonymous bit-fields cannot be cv-qualified (CWG 2229). if (!InvalidDecl && getLangOpts().CPlusPlus && !II && BitWidth && T.hasQualifiers()) { diff --git a/clang/test/CodeGenHIP/amdgpu-barrier-type.hip b/clang/test/CodeGenHIP/amdgpu-barrier-type.hip index 947ceb56d279e..d32e2dbee2238 100644 --- a/clang/test/CodeGenHIP/amdgpu-barrier-type.hip +++ b/clang/test/CodeGenHIP/amdgpu-barrier-type.hip @@ -1,18 +1,20 @@ -// 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 +// 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 { - __amdgpu_named_workgroup_barrier_t x; - __amdgpu_named_workgroup_barrier_t y; -} str; -__amdgpu_named_workgroup_barrier_t *getBar(); -void useBar(__amdgpu_named_workgroup_barrier_t *); +//. +// CHECK: @bar = addrspace(3) global target("amdgcn.named.barrier", 0) undef, align 4 +// CHECK: @arr = addrspace(3) global [2 x target("amdgcn.named.barrier", 0)] undef, align 4 +// CHECK: @__hip_cuid_ = addrspace(1) global i8 0 +// CHECK: @llvm.compiler.used = appending addrspace(1) global [1 x ptr] [ptr addrspacecast (ptr addrspace(1) @__hip_cuid_ to ptr)], section "llvm.metadata" +//. +__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 +24,26 @@ 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 getelementptr inbounds nuw (i8, ptr addrspacecast (ptr addrspace(3) @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(getBar()); return getBar(); } +//. +// CHECK: attributes #[[ATTR0]] = { convergent mustprogress noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="gfx1250" "uniform-work-group-size" } +// CHECK: attributes #[[ATTR1:[0-9]+]] = { convergent nounwind "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="gfx1250" "uniform-work-group-size" } +// CHECK: attributes #[[ATTR2]] = { convergent nounwind "uniform-work-group-size" } +//. +// CHECK: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 600} +// CHECK: [[META1:![0-9]+]] = !{i32 1, !"amdgpu_printf_kind", !"hostcall"} +// CHECK: [[META2:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"} +//. diff --git a/clang/test/SemaCXX/amdgpu-barrier.cpp b/clang/test/SemaCXX/amdgpu-barrier.cpp index a171433727dda..497fb96503c33 100644 --- a/clang/test/SemaCXX/amdgpu-barrier.cpp +++ b/clang/test/SemaCXX/amdgpu-barrier.cpp @@ -13,5 +13,9 @@ void foo() { void *vp = (void *)k; // expected-error {{cannot cast from type '__amdgpu_named_workgroup_barrier_t' to pointer type 'void *'}} } +struct { + __amdgpu_named_workgroup_barrier_t x; // expected-error {{the '__amdgpu_named_workgroup_barrier_t' type cannot be used to declare a structure or union field}} +} str; + 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..4213ea540faf6 100644 --- a/clang/test/SemaHIP/amdgpu-barrier.hip +++ b/clang/test/SemaHIP/amdgpu-barrier.hip @@ -16,5 +16,10 @@ __device__ void foo() { void *vp = (void *)k; // expected-error {{cannot cast from type '__amdgpu_named_workgroup_barrier_t' to pointer type 'void *'}} } +struct { + __amdgpu_named_workgroup_barrier_t x; // expected-error {{the '__amdgpu_named_workgroup_barrier_t' type cannot be used to declare a structure or union field}} + __amdgpu_named_workgroup_barrier_t y[2]; // expected-error {{the '__amdgpu_named_workgroup_barrier_t[2]' type cannot be used to declare a structure or union field}} +} str; + 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..1e2d32005c444 100644 --- a/clang/test/SemaOpenCL/amdgpu-barrier.cl +++ b/clang/test/SemaOpenCL/amdgpu-barrier.cl @@ -3,6 +3,11 @@ // RUN: %clang_cc1 -verify -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -Wno-unused-value %s void foo() { + struct { + __amdgpu_named_workgroup_barrier_t x; // expected-error {{the '__amdgpu_named_workgroup_barrier_t' type cannot be used to declare a structure or union field}} + __amdgpu_named_workgroup_barrier_t y[2]; // expected-error {{the '__amdgpu_named_workgroup_barrier_t[2]' type cannot be used to declare a structure or union field}} + } str; + 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 1de6f77f031e809c14d9b35d763a7688c6cfca44 Mon Sep 17 00:00:00 2001 From: pvanhout <[email protected]> Date: Mon, 6 Jul 2026 13:56:08 +0200 Subject: [PATCH 02/10] Address Comments --- clang/include/clang/Basic/AddressSpaces.h | 2 +- clang/lib/AST/Type.cpp | 7 ++++--- clang/lib/AST/TypePrinter.cpp | 4 ++-- clang/lib/Basic/TargetInfo.cpp | 2 +- clang/lib/Basic/Targets/AArch64.h | 2 +- clang/lib/Basic/Targets/AMDGPU.cpp | 2 +- clang/lib/Basic/Targets/DirectX.h | 2 +- clang/lib/Basic/Targets/NVPTX.h | 2 +- clang/lib/Basic/Targets/SPIR.h | 4 ++-- clang/lib/Basic/Targets/SystemZ.h | 2 +- clang/lib/Basic/Targets/TCE.h | 2 +- clang/lib/Basic/Targets/WebAssembly.h | 2 +- clang/lib/Basic/Targets/X86.h | 2 +- clang/lib/CodeGen/CodeGenModule.cpp | 5 +---- clang/test/SemaCXX/amdgpu-barrier.cpp | 4 ++++ clang/test/SemaHIP/amdgpu-barrier.hip | 3 +++ clang/test/SemaOpenCL/amdgpu-barrier.cl | 5 +++-- 17 files changed, 29 insertions(+), 23 deletions(-) diff --git a/clang/include/clang/Basic/AddressSpaces.h b/clang/include/clang/Basic/AddressSpaces.h index d16fc8069ae43..339fc493bfbaf 100644 --- a/clang/include/clang/Basic/AddressSpaces.h +++ b/clang/include/clang/Basic/AddressSpaces.h @@ -69,7 +69,7 @@ enum class LangAS : unsigned { wasm_funcref, // HIP-specific address spaces - hip_barrier, + 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 diff --git a/clang/lib/AST/Type.cpp b/clang/lib/AST/Type.cpp index cb992eac05518..97816ded541c8 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::hip_barrier)) || + 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 @@ -5493,9 +5493,10 @@ bool Type::isCUDADeviceBuiltinTextureType() const { } bool Type::isAMDGPUNamedBarrierType() const { - const Type *Ty = getUnqualifiedDesugaredType(); + // This query does not care about qualifiers at all. + const Type *Ty = getCanonicalTypeInternal().getTypePtr(); - // unwrap arrays + // Unwrap arrays. while (isa<ArrayType>(Ty)) Ty = Ty->getArrayElementTypeNoTypeQual(); diff --git a/clang/lib/AST/TypePrinter.cpp b/clang/lib/AST/TypePrinter.cpp index e467397d9df52..aba901e764084 100644 --- a/clang/lib/AST/TypePrinter.cpp +++ b/clang/lib/AST/TypePrinter.cpp @@ -2748,8 +2748,8 @@ std::string Qualifiers::getAddrSpaceAsString(LangAS AS) { return "hlsl_push_constant"; case LangAS::wasm_funcref: return "__funcref"; - case LangAS::hip_barrier: - return "hip_barrier"; + 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 95cbe66ec54d5..e14e5f7489168 100644 --- a/clang/lib/Basic/TargetInfo.cpp +++ b/clang/lib/Basic/TargetInfo.cpp @@ -55,7 +55,7 @@ static const LangASMap FakeAddrSpaceMap = { 18, // hlsl_output 19, // hlsl_push_constant 20, // wasm_funcref - 21, // hip_barrier + 21, // amdgpu_barrier }; // TargetInfo Constructor. diff --git a/clang/lib/Basic/Targets/AArch64.h b/clang/lib/Basic/Targets/AArch64.h index ccf5333370925..f02d17fa3cbac 100644 --- a/clang/lib/Basic/Targets/AArch64.h +++ b/clang/lib/Basic/Targets/AArch64.h @@ -54,7 +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, // hip_barrier + 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 1619d2bbfb8b7..f9855161c0837 100644 --- a/clang/lib/Basic/Targets/AMDGPU.cpp +++ b/clang/lib/Basic/Targets/AMDGPU.cpp @@ -57,7 +57,7 @@ const LangASMap AMDGPUTargetInfo::AMDGPUAddrSpaceMap = { llvm::AMDGPUAS::PRIVATE_ADDRESS, // hlsl_output llvm::AMDGPUAS::GLOBAL_ADDRESS, // hlsl_push_constant llvm::AMDGPUAS::FLAT_ADDRESS, // wasm_funcref - llvm::AMDGPUAS::LOCAL_ADDRESS, // hip_barrier + llvm::AMDGPUAS::LOCAL_ADDRESS, // amdgpu_barrier }; } // namespace targets diff --git a/clang/lib/Basic/Targets/DirectX.h b/clang/lib/Basic/Targets/DirectX.h index 1b6785d92d3e3..67725141eb47c 100644 --- a/clang/lib/Basic/Targets/DirectX.h +++ b/clang/lib/Basic/Targets/DirectX.h @@ -51,7 +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, // hip_barrier + 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 534d88e5e333a..803629e3406d4 100644 --- a/clang/lib/Basic/Targets/NVPTX.h +++ b/clang/lib/Basic/Targets/NVPTX.h @@ -55,7 +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, // hip_barrier + 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 ae2b598b46d67..0944dccb04e4a 100644 --- a/clang/lib/Basic/Targets/SPIR.h +++ b/clang/lib/Basic/Targets/SPIR.h @@ -59,7 +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, // hip_barrier + 0, // amdgpu_barrier }; // Used by both the SPIR and SPIR-V targets. @@ -98,7 +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, // hip_barrier + 0, // amdgpu_barrier }; // Base class for SPIR and SPIR-V target info. diff --git a/clang/lib/Basic/Targets/SystemZ.h b/clang/lib/Basic/Targets/SystemZ.h index 6d1420afc58ca..c12cefc10e08a 100644 --- a/clang/lib/Basic/Targets/SystemZ.h +++ b/clang/lib/Basic/Targets/SystemZ.h @@ -49,7 +49,7 @@ static const unsigned ZOSAddressMap[] = { 0, // hlsl_output 0, // hlsl_push_constant 0, // wasm_funcref - 0, // hip_barrier + 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 95a17e0a9356f..403b8b1bb3304 100644 --- a/clang/lib/Basic/Targets/TCE.h +++ b/clang/lib/Basic/Targets/TCE.h @@ -60,7 +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, // hip_barrier + 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 e991e1bd19d6f..5389025817b02 100644 --- a/clang/lib/Basic/Targets/WebAssembly.h +++ b/clang/lib/Basic/Targets/WebAssembly.h @@ -49,7 +49,7 @@ static const unsigned WebAssemblyAddrSpaceMap[] = { 0, // hlsl_output 0, // hlsl_push_constant 20, // wasm_funcref - 0, // hip_barrier + 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 6a52c80662d3b..ab5cd9618bbaa 100644 --- a/clang/lib/Basic/Targets/X86.h +++ b/clang/lib/Basic/Targets/X86.h @@ -55,7 +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, // hip_barrier + 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 ed76ea5c53b8a..9aed04bc769c3 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -6227,11 +6227,8 @@ LangAS CodeGenModule::GetGlobalVarAddressSpace(const VarDecl *D) { if (LangOpts.CUDA && LangOpts.CUDAIsDevice) { if (D) { - // TOOD: Forbid type on struct fields - // llvm::dbgs() << D->getNameAsString() << ": " << - // D->getType()->isAMDGPUNamedBarrierType() << "\n"; if (D->getType()->isAMDGPUNamedBarrierType()) - return LangAS::hip_barrier; + return LangAS::amdgpu_barrier; if (D->hasAttr<CUDAConstantAttr>()) return LangAS::cuda_constant; diff --git a/clang/test/SemaCXX/amdgpu-barrier.cpp b/clang/test/SemaCXX/amdgpu-barrier.cpp index 497fb96503c33..ff0cd5432c9e7 100644 --- a/clang/test/SemaCXX/amdgpu-barrier.cpp +++ b/clang/test/SemaCXX/amdgpu-barrier.cpp @@ -13,8 +13,12 @@ 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 { __amdgpu_named_workgroup_barrier_t x; // expected-error {{the '__amdgpu_named_workgroup_barrier_t' type cannot be used to declare a structure or union field}} + __amdgpu_named_workgroup_barrier_t y[2]; // expected-error {{the '__amdgpu_named_workgroup_barrier_t[2]' type cannot be used to declare a structure or union field}} + SugaredArray z[2]; // expected-error {{the 'SugaredArray[2]' (aka '__amdgpu_named_workgroup_barrier_t[2][2]') type cannot be used to declare a structure or union field}} } str; static_assert(sizeof(__amdgpu_named_workgroup_barrier_t) == 16, "wrong size"); diff --git a/clang/test/SemaHIP/amdgpu-barrier.hip b/clang/test/SemaHIP/amdgpu-barrier.hip index 4213ea540faf6..b30ddb3845595 100644 --- a/clang/test/SemaHIP/amdgpu-barrier.hip +++ b/clang/test/SemaHIP/amdgpu-barrier.hip @@ -16,9 +16,12 @@ __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 { __amdgpu_named_workgroup_barrier_t x; // expected-error {{the '__amdgpu_named_workgroup_barrier_t' type cannot be used to declare a structure or union field}} __amdgpu_named_workgroup_barrier_t y[2]; // expected-error {{the '__amdgpu_named_workgroup_barrier_t[2]' type cannot be used to declare a structure or union field}} + SugaredArray z[2]; // expected-error {{the 'SugaredArray[2]' (aka '__amdgpu_named_workgroup_barrier_t[2][2]') type cannot be used to declare a structure or union field}} } str; static_assert(sizeof(__amdgpu_named_workgroup_barrier_t) == 16, "wrong size"); diff --git a/clang/test/SemaOpenCL/amdgpu-barrier.cl b/clang/test/SemaOpenCL/amdgpu-barrier.cl index 1e2d32005c444..1963ac1a120f1 100644 --- a/clang/test/SemaOpenCL/amdgpu-barrier.cl +++ b/clang/test/SemaOpenCL/amdgpu-barrier.cl @@ -4,8 +4,9 @@ void foo() { struct { - __amdgpu_named_workgroup_barrier_t x; // expected-error {{the '__amdgpu_named_workgroup_barrier_t' type cannot be used to declare a structure or union field}} - __amdgpu_named_workgroup_barrier_t y[2]; // expected-error {{the '__amdgpu_named_workgroup_barrier_t[2]' type cannot be used to declare a structure or union field}} + __amdgpu_named_workgroup_barrier_t x; // expected-error {{the '__amdgpu_named_workgroup_barrier_t' type cannot be used to declare a structure or union field}} + __amdgpu_named_workgroup_barrier_t y[2]; // expected-error {{the '__amdgpu_named_workgroup_barrier_t[2]' type cannot be used to declare a structure or union field}} + __amdgpu_named_workgroup_barrier_t z[2][2]; // expected-error {{the '__amdgpu_named_workgroup_barrier_t[2][2]' type cannot be used to declare a structure or union field}} } str; int n = 100; >From 72c6559899a6bfe5214ec23badeffebdff9581ea Mon Sep 17 00:00:00 2001 From: pvanhout <[email protected]> Date: Tue, 7 Jul 2026 09:35:57 +0200 Subject: [PATCH 03/10] Re-allow wrappers around named barriers --- clang/include/clang/AST/TypeBase.h | 5 +- clang/include/clang/Basic/Attr.td | 9 ++- .../clang/Basic/DiagnosticSemaKinds.td | 11 ++++ clang/include/clang/Sema/SemaAMDGPU.h | 3 + clang/lib/AST/Type.cpp | 20 +++++-- clang/lib/CodeGen/CodeGenModule.cpp | 2 +- clang/lib/Sema/SemaAMDGPU.cpp | 59 +++++++++++++++++++ clang/lib/Sema/SemaDecl.cpp | 12 ++-- clang/test/CodeGenHIP/amdgpu-barrier-type.hip | 8 ++- clang/test/SemaCXX/amdgpu-barrier.cpp | 35 +++++++++-- clang/test/SemaHIP/amdgpu-barrier.hip | 35 +++++++++-- clang/test/SemaOpenCL/amdgpu-barrier.cl | 28 +++++++-- 12 files changed, 196 insertions(+), 31 deletions(-) diff --git a/clang/include/clang/AST/TypeBase.h b/clang/include/clang/AST/TypeBase.h index 2a3aa9ebe4887..2182cd79144fa 100644 --- a/clang/include/clang/AST/TypeBase.h +++ b/clang/include/clang/AST/TypeBase.h @@ -2813,8 +2813,11 @@ 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. + /// 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/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 e378fe8750718..f63000e71c94e 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -14295,6 +14295,17 @@ 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< + "%0 is only allowed as a field if the %select{struct|interface|union|class|enum}1" + " is a trivial wrapper around it">; +def note_amdgcn_not_a_wrapper_derived_class : Note< + "%0 is not a trivial wrapper because it inherits from %1">; +def note_amdgcn_not_a_wrapper_too_many_fields : Note< + "%0 is not a trivial wrapper because it has more than one field">; +def note_amdgcn_wrapper_not_final : Note< + "%0 is not a trivial wrapper because it is not marked 'final'">; + // 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/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/Type.cpp b/clang/lib/AST/Type.cpp index 97816ded541c8..e51e7de9f176a 100644 --- a/clang/lib/AST/Type.cpp +++ b/clang/lib/AST/Type.cpp @@ -5492,19 +5492,31 @@ bool Type::isCUDADeviceBuiltinTextureType() const { return false; } -bool Type::isAMDGPUNamedBarrierType() const { +static bool isAMDGPUNamedBarrierTypeImpl(const Type *Ty, bool AllowWrappers) { // This query does not care about qualifiers at all. - const Type *Ty = getCanonicalTypeInternal().getTypePtr(); + Ty = Ty->getUnqualifiedDesugaredType(); // Unwrap arrays. while (isa<ArrayType>(Ty)) - Ty = Ty->getArrayElementTypeNoTypeQual(); + Ty = Ty->getArrayElementTypeNoTypeQual()->getUnqualifiedDesugaredType(); - if (const auto *BT = dyn_cast<BuiltinType>(Ty->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/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 9aed04bc769c3..ba2e6a13bcf69 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -6227,7 +6227,7 @@ LangAS CodeGenModule::GetGlobalVarAddressSpace(const VarDecl *D) { if (LangOpts.CUDA && LangOpts.CUDAIsDevice) { if (D) { - if (D->getType()->isAMDGPUNamedBarrierType()) + if (D->getType()->isAMDGPUNamedBarrierTypeOrWrapper()) return LangAS::amdgpu_barrier; if (D->hasAttr<CUDAConstantAttr>()) diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp index f1de44e3d2ed7..b41912c826e5c 100644 --- a/clang/lib/Sema/SemaAMDGPU.cpp +++ b/clang/lib/Sema/SemaAMDGPU.cpp @@ -1062,4 +1062,63 @@ bool DiagnoseUnguardedBuiltins::VisitCallExpr(CallExpr *CE) { void SemaAMDGPU::DiagnoseUnguardedBuiltinUsage(FunctionDecl *FD) { DiagnoseUnguardedBuiltins(SemaRef).IssueDiagnostics(FD->getBody()); } + +void SemaAMDGPU::checkNamedBarrierWrapper(RecordDecl *R) { + if (R->isInvalidDecl()) + return; + + // Check if this record has any fields that interest us. + FieldDecl *NamedBarrField = nullptr; + for (FieldDecl *FD : R->fields()) { + QualType FDTy = FD->getType(); + if (FDTy->isAMDGPUNamedBarrierType()) { + NamedBarrField = FD; + break; + } + } + + if (!NamedBarrField) + return; + + bool IsInvalid = false; + const auto OnError = [&]() { + if (!IsInvalid) { + SemaRef.Diag(NamedBarrField->getLocation(), + diag::err_amdgcn_invalid_field_not_a_wrapper) + << NamedBarrField->getType() << R->getTagKind(); + IsInvalid = true; + } + }; + + if (R->getNumFields() > 1) { + OnError(); + SemaRef.Diag(R->getLocation(), + diag::note_amdgcn_not_a_wrapper_too_many_fields) + << R->getName(); + } + + if (const auto *CxxR = dyn_cast<CXXRecordDecl>(R)) { + if (CxxR->getNumBases() != 0) { + OnError(); + for (CXXBaseSpecifier CxxBase : CxxR->bases()) { + SemaRef.Diag(CxxBase.getBaseTypeLoc(), + diag::note_amdgcn_not_a_wrapper_derived_class) + << R->getName() << CxxBase.getType(); + } + } + + if (!CxxR->hasAttr<FinalAttr>()) { + OnError(); + SemaRef.Diag(R->getLocation(), diag::note_amdgcn_wrapper_not_final) + << R->getName(); + } + } + + if (IsInvalid) + return; + + ASTContext &Context = getASTContext(); + R->addAttr(AMDGPUNamedBarrierWrapperAttr::CreateImplicit( + Context, NamedBarrField->getSourceRange())); +} } // namespace clang diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 2dd524611423a..635dd2d30e578 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -19489,14 +19489,6 @@ FieldDecl *Sema::CheckFieldDecl(DeclarationName Name, QualType T, } } - // AMDGPU does not allows the following types to be used for structure or - // union fields: named barriers - if (T->isAMDGPUNamedBarrierType()) { - Diag(Loc, diag::err_invalid_type_for_struct_or_union_field) << T; - Record->setInvalidDecl(); - InvalidDecl = true; - } - // Anonymous bit-fields cannot be cv-qualified (CWG 2229). if (!InvalidDecl && getLangOpts().CPlusPlus && !II && BitWidth && T.hasQualifiers()) { @@ -20474,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 d32e2dbee2238..724136597810a 100644 --- a/clang/test/CodeGenHIP/amdgpu-barrier-type.hip +++ b/clang/test/CodeGenHIP/amdgpu-barrier-type.hip @@ -6,13 +6,17 @@ __shared__ __amdgpu_named_workgroup_barrier_t bar; __shared__ __amdgpu_named_workgroup_barrier_t arr[2]; - //. // CHECK: @bar = addrspace(3) global target("amdgcn.named.barrier", 0) undef, align 4 // CHECK: @arr = addrspace(3) global [2 x target("amdgcn.named.barrier", 0)] undef, align 4 +// CHECK: @wrapper_str = addrspace(3) global %struct.WrapperStruct undef, align 4 // CHECK: @__hip_cuid_ = addrspace(1) global i8 0 // CHECK: @llvm.compiler.used = appending addrspace(1) global [1 x ptr] [ptr addrspacecast (ptr addrspace(1) @__hip_cuid_ to ptr)], section "llvm.metadata" //. +__shared__ struct WrapperStruct final { + __amdgpu_named_workgroup_barrier_t x; +} wrapper_str; + __attribute__((device)) __amdgpu_named_workgroup_barrier_t *getBar(); __attribute__((device)) void useBar(__amdgpu_named_workgroup_barrier_t *); @@ -25,6 +29,7 @@ __attribute__((device)) void useBar(__amdgpu_named_workgroup_barrier_t *); // 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(3) @bar to ptr)) #[[ATTR2]] +// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef addrspacecast (ptr addrspace(3) @wrapper_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) @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]] @@ -34,6 +39,7 @@ __attribute__((device)) void useBar(__amdgpu_named_workgroup_barrier_t *); __attribute__((device)) __amdgpu_named_workgroup_barrier_t *testSem(__amdgpu_named_workgroup_barrier_t *p) { useBar(p); useBar(&bar); + useBar(&wrapper_str.x); useBar(&arr[1]); useBar(getBar()); return getBar(); diff --git a/clang/test/SemaCXX/amdgpu-barrier.cpp b/clang/test/SemaCXX/amdgpu-barrier.cpp index ff0cd5432c9e7..c94663327d75c 100644 --- a/clang/test/SemaCXX/amdgpu-barrier.cpp +++ b/clang/test/SemaCXX/amdgpu-barrier.cpp @@ -15,11 +15,36 @@ void foo() { using SugaredArray = __amdgpu_named_workgroup_barrier_t[2]; -struct { - __amdgpu_named_workgroup_barrier_t x; // expected-error {{the '__amdgpu_named_workgroup_barrier_t' type cannot be used to declare a structure or union field}} - __amdgpu_named_workgroup_barrier_t y[2]; // expected-error {{the '__amdgpu_named_workgroup_barrier_t[2]' type cannot be used to declare a structure or union field}} - SugaredArray z[2]; // expected-error {{the 'SugaredArray[2]' (aka '__amdgpu_named_workgroup_barrier_t[2][2]') type cannot be used to declare a structure or union field}} -} str; +struct TestSimple final { + __amdgpu_named_workgroup_barrier_t x; +}; + +struct TestArray final{ + __amdgpu_named_workgroup_barrier_t y[2]; +}; + +struct TestSugared final { + SugaredArray z[2]; +}; + +// Wrappers cannot have >1 field. +struct WrapperHasTooManyFields final { // expected-note {{WrapperHasTooManyFields is not a trivial wrapper because it has more than one field}} + __amdgpu_named_workgroup_barrier_t x; // expected-error {{'__amdgpu_named_workgroup_barrier_t' is only allowed as a field if the struct is a trivial wrapper around it}} + int other; +}; + +struct WrapperIsNotFinal { // expected-note {{WrapperIsNotFinal is not a trivial wrapper because it is not marked 'final'}} + __amdgpu_named_workgroup_barrier_t x; // expected-error {{'__amdgpu_named_workgroup_barrier_t' is only allowed as a field if the struct is a trivial wrapper around it}} +}; + +// Wrappers cannot have base class +struct WrapperBase { +}; + +struct WrapperWithBase final : public WrapperBase { // expected-note {{WrapperWithBase is not a trivial wrapper because it inherits from 'WrapperBase'}} + __amdgpu_named_workgroup_barrier_t x; // expected-error {{'__amdgpu_named_workgroup_barrier_t' is only allowed as a field if the struct is a trivial wrapper around it}} +}; + 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 b30ddb3845595..679d2fbcf980f 100644 --- a/clang/test/SemaHIP/amdgpu-barrier.hip +++ b/clang/test/SemaHIP/amdgpu-barrier.hip @@ -18,11 +18,36 @@ __device__ void foo() { using SugaredArray = __amdgpu_named_workgroup_barrier_t[2]; -struct { - __amdgpu_named_workgroup_barrier_t x; // expected-error {{the '__amdgpu_named_workgroup_barrier_t' type cannot be used to declare a structure or union field}} - __amdgpu_named_workgroup_barrier_t y[2]; // expected-error {{the '__amdgpu_named_workgroup_barrier_t[2]' type cannot be used to declare a structure or union field}} - SugaredArray z[2]; // expected-error {{the 'SugaredArray[2]' (aka '__amdgpu_named_workgroup_barrier_t[2][2]') type cannot be used to declare a structure or union field}} -} str; +struct TestSimple final { + __amdgpu_named_workgroup_barrier_t x; +}; + +struct TestArray final { + __amdgpu_named_workgroup_barrier_t y[2]; +}; + +struct TestSugared final { + SugaredArray z[2]; +}; + +// Wrappers cannot have >1 field. +struct WrapperHasTooManyFields final { // expected-note {{WrapperHasTooManyFields is not a trivial wrapper because it has more than one field}} + __amdgpu_named_workgroup_barrier_t x; // expected-error {{'__amdgpu_named_workgroup_barrier_t' is only allowed as a field if the struct is a trivial wrapper around it}} + int other; +}; + +struct WrapperIsNotFinal { // expected-note {{WrapperIsNotFinal is not a trivial wrapper because it is not marked 'final'}} + __amdgpu_named_workgroup_barrier_t x; // expected-error {{'__amdgpu_named_workgroup_barrier_t' is only allowed as a field if the struct is a trivial wrapper around it}} +}; + +// Wrappers cannot have base class +struct WrapperBase { +}; + +struct WrapperWithBase final : public WrapperBase { // expected-note {{WrapperWithBase is not a trivial wrapper because it inherits from 'WrapperBase'}} + __amdgpu_named_workgroup_barrier_t x; // expected-error {{'__amdgpu_named_workgroup_barrier_t' is only allowed as a field if the struct is a trivial wrapper around it}} +}; + 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 1963ac1a120f1..2c6a305780496 100644 --- a/clang/test/SemaOpenCL/amdgpu-barrier.cl +++ b/clang/test/SemaOpenCL/amdgpu-barrier.cl @@ -3,11 +3,29 @@ // RUN: %clang_cc1 -verify -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -Wno-unused-value %s void foo() { - struct { - __amdgpu_named_workgroup_barrier_t x; // expected-error {{the '__amdgpu_named_workgroup_barrier_t' type cannot be used to declare a structure or union field}} - __amdgpu_named_workgroup_barrier_t y[2]; // expected-error {{the '__amdgpu_named_workgroup_barrier_t[2]' type cannot be used to declare a structure or union field}} - __amdgpu_named_workgroup_barrier_t z[2][2]; // expected-error {{the '__amdgpu_named_workgroup_barrier_t[2][2]' type cannot be used to declare a structure or union field}} - } str; + 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 trivial wrapper because it has more than one field}} + __amdgpu_named_workgroup_barrier_t x; // expected-error {{'__amdgpu_named_workgroup_barrier_t' is only allowed as a field if the struct is a trivial wrapper around it}} + 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'}} >From ed47c3537766efd78c67173e1c557973c61ddca0 Mon Sep 17 00:00:00 2001 From: pvanhout <[email protected]> Date: Tue, 7 Jul 2026 09:42:12 +0200 Subject: [PATCH 04/10] Add test w/ static field --- clang/test/SemaCXX/amdgpu-barrier.cpp | 5 +++++ clang/test/SemaHIP/amdgpu-barrier.hip | 5 +++++ 2 files changed, 10 insertions(+) diff --git a/clang/test/SemaCXX/amdgpu-barrier.cpp b/clang/test/SemaCXX/amdgpu-barrier.cpp index c94663327d75c..0b054d9d2b0e0 100644 --- a/clang/test/SemaCXX/amdgpu-barrier.cpp +++ b/clang/test/SemaCXX/amdgpu-barrier.cpp @@ -27,6 +27,11 @@ struct TestSugared final { SugaredArray z[2]; }; +struct TestSimpleWithStaticField final { + __amdgpu_named_workgroup_barrier_t x; + static unsigned Harmless; +}; + // Wrappers cannot have >1 field. struct WrapperHasTooManyFields final { // expected-note {{WrapperHasTooManyFields is not a trivial wrapper because it has more than one field}} __amdgpu_named_workgroup_barrier_t x; // expected-error {{'__amdgpu_named_workgroup_barrier_t' is only allowed as a field if the struct is a trivial wrapper around it}} diff --git a/clang/test/SemaHIP/amdgpu-barrier.hip b/clang/test/SemaHIP/amdgpu-barrier.hip index 679d2fbcf980f..5378fdc6793d0 100644 --- a/clang/test/SemaHIP/amdgpu-barrier.hip +++ b/clang/test/SemaHIP/amdgpu-barrier.hip @@ -30,6 +30,11 @@ struct TestSugared final { SugaredArray z[2]; }; +struct TestSimpleWithStaticField final { + __amdgpu_named_workgroup_barrier_t x; + static unsigned Harmless; +}; + // Wrappers cannot have >1 field. struct WrapperHasTooManyFields final { // expected-note {{WrapperHasTooManyFields is not a trivial wrapper because it has more than one field}} __amdgpu_named_workgroup_barrier_t x; // expected-error {{'__amdgpu_named_workgroup_barrier_t' is only allowed as a field if the struct is a trivial wrapper around it}} >From 67682530faaa16a4095d10d5bae86382e2056c1a Mon Sep 17 00:00:00 2001 From: pvanhout <[email protected]> Date: Wed, 8 Jul 2026 11:40:54 +0200 Subject: [PATCH 05/10] Update implementation --- clang/include/clang/Basic/Attr.td | 2 +- .../clang/Basic/DiagnosticSemaKinds.td | 15 ++++----- clang/lib/Sema/SemaAMDGPU.cpp | 25 +++++++-------- clang/lib/Sema/SemaDeclCXX.cpp | 13 ++++++++ clang/test/CodeGenHIP/amdgpu-barrier-type.hip | 10 +++++- clang/test/SemaCXX/amdgpu-barrier.cpp | 32 +++++++++++-------- clang/test/SemaHIP/amdgpu-barrier.hip | 32 +++++++++++-------- clang/test/SemaOpenCL/amdgpu-barrier.cl | 4 +-- 8 files changed, 82 insertions(+), 51 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index f486bd176c9f8..e81498b0349c0 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -2534,7 +2534,7 @@ def AMDGPUMaxNumWorkGroups : InheritableAttr { def AMDGPUNamedBarrierWrapper : InheritableAttr { let Spellings = []; - let Args = []; + let Args = [DeclArgument<Field, "WrappedField">]; let Documentation = [InternalOnly]; let SemaHandler = 0; } diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index f63000e71c94e..3ea0ed78fb3d3 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -14297,14 +14297,13 @@ def note_acc_reduction_combiner_forming // AMDGCN type diagnostics def err_amdgcn_invalid_field_not_a_wrapper : Error< - "%0 is only allowed as a field if the %select{struct|interface|union|class|enum}1" - " is a trivial wrapper around it">; -def note_amdgcn_not_a_wrapper_derived_class : Note< - "%0 is not a trivial wrapper because it inherits from %1">; -def note_amdgcn_not_a_wrapper_too_many_fields : Note< - "%0 is not a trivial wrapper because it has more than one field">; -def note_amdgcn_wrapper_not_final : Note< - "%0 is not a trivial wrapper because it is not marked 'final'">; + "fields of type %0 are only allowed in named barrier wrappers">; +def note_amdgcn_not_a_named_barrier_wrapper_derived_class : Note< + "%0 is not a named barrier wrapper because it inherits from %1; named barrier wrappers may not have any bases">; +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 note_amdgcn_named_barrier_wrapper_implies_final : Note< + "%0 is a named barrier wrapper around %1, which implies 'final'">; // AMDGCN builtins diagnostics def err_amdgcn_load_lds_size_invalid_value : Error<"invalid size value">; diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp index b41912c826e5c..b951c2f6552dd 100644 --- a/clang/lib/Sema/SemaAMDGPU.cpp +++ b/clang/lib/Sema/SemaAMDGPU.cpp @@ -1068,10 +1068,12 @@ void SemaAMDGPU::checkNamedBarrierWrapper(RecordDecl *R) { return; // Check if this record has any fields that interest us. + // We cannot codegen any kind of heterogenous struct where we named barrier + // fields mixed with other fields. FieldDecl *NamedBarrField = nullptr; for (FieldDecl *FD : R->fields()) { QualType FDTy = FD->getType(); - if (FDTy->isAMDGPUNamedBarrierType()) { + if (FDTy->isAMDGPUNamedBarrierTypeOrWrapper()) { NamedBarrField = FD; break; } @@ -1085,7 +1087,7 @@ void SemaAMDGPU::checkNamedBarrierWrapper(RecordDecl *R) { if (!IsInvalid) { SemaRef.Diag(NamedBarrField->getLocation(), diag::err_amdgcn_invalid_field_not_a_wrapper) - << NamedBarrField->getType() << R->getTagKind(); + << NamedBarrField->getType(); IsInvalid = true; } }; @@ -1093,7 +1095,7 @@ void SemaAMDGPU::checkNamedBarrierWrapper(RecordDecl *R) { if (R->getNumFields() > 1) { OnError(); SemaRef.Diag(R->getLocation(), - diag::note_amdgcn_not_a_wrapper_too_many_fields) + diag::note_amdgcn_not_a_named_barrier_wrapper_too_many_fields) << R->getName(); } @@ -1101,24 +1103,21 @@ void SemaAMDGPU::checkNamedBarrierWrapper(RecordDecl *R) { if (CxxR->getNumBases() != 0) { OnError(); for (CXXBaseSpecifier CxxBase : CxxR->bases()) { - SemaRef.Diag(CxxBase.getBaseTypeLoc(), - diag::note_amdgcn_not_a_wrapper_derived_class) + SemaRef.Diag( + CxxBase.getBaseTypeLoc(), + diag::note_amdgcn_not_a_named_barrier_wrapper_derived_class) << R->getName() << CxxBase.getType(); } } - - if (!CxxR->hasAttr<FinalAttr>()) { - OnError(); - SemaRef.Diag(R->getLocation(), diag::note_amdgcn_wrapper_not_final) - << R->getName(); - } } if (IsInvalid) return; ASTContext &Context = getASTContext(); - R->addAttr(AMDGPUNamedBarrierWrapperAttr::CreateImplicit( - Context, NamedBarrField->getSourceRange())); + SourceRange SR = NamedBarrField->getSourceRange(); + R->addAttr(FinalAttr::CreateImplicit(Context, SR)); + R->addAttr(AMDGPUNamedBarrierWrapperAttr::CreateImplicit(Context, + NamedBarrField, SR)); } } // namespace clang diff --git a/clang/lib/Sema/SemaDeclCXX.cpp b/clang/lib/Sema/SemaDeclCXX.cpp index 685412d681680..411e247523b73 100644 --- a/clang/lib/Sema/SemaDeclCXX.cpp +++ b/clang/lib/Sema/SemaDeclCXX.cpp @@ -2854,6 +2854,19 @@ CXXBaseSpecifier *Sema::CheckBaseSpecifier(CXXRecordDecl *Class, << BaseDecl->getDeclName() << FA->isSpelledAsSealed(); Diag(BaseDecl->getLocation(), diag::note_entity_declared_at) << BaseDecl->getDeclName() << FA->getRange(); + + // AMDGPU Named Barrier Wrappers are final. If the FinalAttr may have come + // from there, diagnose it. + if (FA->isImplicit()) { + if (auto *NBW = BaseDecl->getAttr<AMDGPUNamedBarrierWrapperAttr>()) { + FieldDecl *WrappedField = NBW->getWrappedField(); + Diag(BaseDecl->getLocation(), + diag::note_amdgcn_named_barrier_wrapper_implies_final) + << BaseDecl->getDeclName() << WrappedField->getDeclName(); + Diag(WrappedField->getLocation(), diag::note_entity_declared_at) + << WrappedField->getDeclName(); + } + } return nullptr; } diff --git a/clang/test/CodeGenHIP/amdgpu-barrier-type.hip b/clang/test/CodeGenHIP/amdgpu-barrier-type.hip index 724136597810a..3664821b80541 100644 --- a/clang/test/CodeGenHIP/amdgpu-barrier-type.hip +++ b/clang/test/CodeGenHIP/amdgpu-barrier-type.hip @@ -6,17 +6,23 @@ __shared__ __amdgpu_named_workgroup_barrier_t bar; __shared__ __amdgpu_named_workgroup_barrier_t arr[2]; + //. // CHECK: @bar = addrspace(3) global target("amdgcn.named.barrier", 0) undef, align 4 // CHECK: @arr = addrspace(3) global [2 x target("amdgcn.named.barrier", 0)] undef, align 4 // CHECK: @wrapper_str = addrspace(3) global %struct.WrapperStruct undef, align 4 +// CHECK: @wrapperwrapper_str = addrspace(3) global %struct.WrapperWrapperStruct undef, align 4 // CHECK: @__hip_cuid_ = addrspace(1) global i8 0 // CHECK: @llvm.compiler.used = appending addrspace(1) global [1 x ptr] [ptr addrspacecast (ptr addrspace(1) @__hip_cuid_ to ptr)], section "llvm.metadata" //. -__shared__ struct WrapperStruct final { +__shared__ struct WrapperStruct { __amdgpu_named_workgroup_barrier_t x; } wrapper_str; +__shared__ struct WrapperWrapperStruct { + WrapperStruct x; +} wrapperwrapper_str; + __attribute__((device)) __amdgpu_named_workgroup_barrier_t *getBar(); __attribute__((device)) void useBar(__amdgpu_named_workgroup_barrier_t *); @@ -30,6 +36,7 @@ __attribute__((device)) void useBar(__amdgpu_named_workgroup_barrier_t *); // 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(3) @bar to ptr)) #[[ATTR2]] // CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef addrspacecast (ptr addrspace(3) @wrapper_str to ptr)) #[[ATTR2]] +// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef addrspacecast (ptr addrspace(3) @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) @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]] @@ -40,6 +47,7 @@ __attribute__((device)) __amdgpu_named_workgroup_barrier_t *testSem(__amdgpu_nam useBar(p); useBar(&bar); useBar(&wrapper_str.x); + useBar(&wrapperwrapper_str.x.x); useBar(&arr[1]); useBar(getBar()); return getBar(); diff --git a/clang/test/SemaCXX/amdgpu-barrier.cpp b/clang/test/SemaCXX/amdgpu-barrier.cpp index 0b054d9d2b0e0..f9e5313bcfa26 100644 --- a/clang/test/SemaCXX/amdgpu-barrier.cpp +++ b/clang/test/SemaCXX/amdgpu-barrier.cpp @@ -15,41 +15,47 @@ void foo() { using SugaredArray = __amdgpu_named_workgroup_barrier_t[2]; -struct TestSimple final { - __amdgpu_named_workgroup_barrier_t x; +// expected-note@+2 {{'TestSimple' is a named barrier wrapper around 'x', which implies 'final'}} +// expected-note@+1 {{'TestSimple' declared here}} +struct TestSimple { + __amdgpu_named_workgroup_barrier_t x; // expected-note {{'x' declared here}} }; -struct TestArray final{ +struct TestArray{ __amdgpu_named_workgroup_barrier_t y[2]; }; -struct TestSugared final { +struct TestSugared { SugaredArray z[2]; }; -struct TestSimpleWithStaticField final { +struct TestSimpleWithStaticField { __amdgpu_named_workgroup_barrier_t x; static unsigned Harmless; }; // Wrappers cannot have >1 field. -struct WrapperHasTooManyFields final { // expected-note {{WrapperHasTooManyFields is not a trivial wrapper because it has more than one field}} - __amdgpu_named_workgroup_barrier_t x; // expected-error {{'__amdgpu_named_workgroup_barrier_t' is only allowed as a field if the struct is a trivial wrapper around it}} +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; }; -struct WrapperIsNotFinal { // expected-note {{WrapperIsNotFinal is not a trivial wrapper because it is not marked 'final'}} - __amdgpu_named_workgroup_barrier_t x; // expected-error {{'__amdgpu_named_workgroup_barrier_t' is only allowed as a field if the struct is a trivial wrapper around it}} -}; - // Wrappers cannot have base class struct WrapperBase { }; -struct WrapperWithBase final : public WrapperBase { // expected-note {{WrapperWithBase is not a trivial wrapper because it inherits from 'WrapperBase'}} - __amdgpu_named_workgroup_barrier_t x; // expected-error {{'__amdgpu_named_workgroup_barrier_t' is only allowed as a field if the struct is a trivial wrapper around it}} +struct WrapperWithBase : public WrapperBase { // expected-note {{WrapperWithBase is not a named barrier wrapper because it inherits from 'WrapperBase'; named barrier wrappers may not have any bases}} + __amdgpu_named_workgroup_barrier_t x; // expected-error {{fields of type '__amdgpu_named_workgroup_barrier_t' are only allowed in named barrier wrappers}} }; +// Wrappers imply final +struct NoWrapperInheritance : public TestSimple { }; // expected-error {{base 'TestSimple' is marked 'final'}} + +// Wrappers of Wrappers have the same restrictions. +struct WrapperOfWrapperWithTooManyFields final { // 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; +}; 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 5378fdc6793d0..a0baaaa3bf899 100644 --- a/clang/test/SemaHIP/amdgpu-barrier.hip +++ b/clang/test/SemaHIP/amdgpu-barrier.hip @@ -18,41 +18,47 @@ __device__ void foo() { using SugaredArray = __amdgpu_named_workgroup_barrier_t[2]; -struct TestSimple final { - __amdgpu_named_workgroup_barrier_t x; +// expected-note@+2 {{'TestSimple' is a named barrier wrapper around 'x', which implies 'final'}} +// expected-note@+1 {{'TestSimple' declared here}} +struct TestSimple { + __amdgpu_named_workgroup_barrier_t x; // expected-note {{'x' declared here}} }; -struct TestArray final { +struct TestArray{ __amdgpu_named_workgroup_barrier_t y[2]; }; -struct TestSugared final { +struct TestSugared { SugaredArray z[2]; }; -struct TestSimpleWithStaticField final { +struct TestSimpleWithStaticField { __amdgpu_named_workgroup_barrier_t x; static unsigned Harmless; }; // Wrappers cannot have >1 field. -struct WrapperHasTooManyFields final { // expected-note {{WrapperHasTooManyFields is not a trivial wrapper because it has more than one field}} - __amdgpu_named_workgroup_barrier_t x; // expected-error {{'__amdgpu_named_workgroup_barrier_t' is only allowed as a field if the struct is a trivial wrapper around it}} +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; }; -struct WrapperIsNotFinal { // expected-note {{WrapperIsNotFinal is not a trivial wrapper because it is not marked 'final'}} - __amdgpu_named_workgroup_barrier_t x; // expected-error {{'__amdgpu_named_workgroup_barrier_t' is only allowed as a field if the struct is a trivial wrapper around it}} -}; - // Wrappers cannot have base class struct WrapperBase { }; -struct WrapperWithBase final : public WrapperBase { // expected-note {{WrapperWithBase is not a trivial wrapper because it inherits from 'WrapperBase'}} - __amdgpu_named_workgroup_barrier_t x; // expected-error {{'__amdgpu_named_workgroup_barrier_t' is only allowed as a field if the struct is a trivial wrapper around it}} +struct WrapperWithBase : public WrapperBase { // expected-note {{WrapperWithBase is not a named barrier wrapper because it inherits from 'WrapperBase'; named barrier wrappers may not have any bases}} + __amdgpu_named_workgroup_barrier_t x; // expected-error {{fields of type '__amdgpu_named_workgroup_barrier_t' are only allowed in named barrier wrappers}} }; +// Wrappers imply final +struct NoWrapperInheritance : public TestSimple { }; // expected-error {{base 'TestSimple' is marked 'final'}} + +// Wrappers of Wrappers have the same restrictions. +struct WrapperOfWrapperWithTooManyFields final { // 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; +}; 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 2c6a305780496..2f96a6aba3a3e 100644 --- a/clang/test/SemaOpenCL/amdgpu-barrier.cl +++ b/clang/test/SemaOpenCL/amdgpu-barrier.cl @@ -22,8 +22,8 @@ void foo() { }; // Wrappers cannot have >1 field. - struct WrapperHasTooManyFields { // expected-note {{WrapperHasTooManyFields is not a trivial wrapper because it has more than one field}} - __amdgpu_named_workgroup_barrier_t x; // expected-error {{'__amdgpu_named_workgroup_barrier_t' is only allowed as a field if the struct is a trivial wrapper around it}} + 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; }; >From e33304c3d59c8eddc5b447eb83ff4053414f8a39 Mon Sep 17 00:00:00 2001 From: pvanhout <[email protected]> Date: Thu, 9 Jul 2026 10:55:22 +0200 Subject: [PATCH 06/10] Use C++ standard layout instead --- clang/include/clang/Basic/Attr.td | 2 +- .../clang/Basic/DiagnosticSemaKinds.td | 11 ++- clang/lib/Sema/SemaAMDGPU.cpp | 96 ++++++++++++------- clang/lib/Sema/SemaDeclCXX.cpp | 13 --- clang/test/SemaCXX/amdgpu-barrier.cpp | 47 +++++++-- clang/test/SemaHIP/amdgpu-barrier.hip | 48 ++++++++-- clang/test/SemaOpenCL/amdgpu-barrier.cl | 2 +- 7 files changed, 143 insertions(+), 76 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index e81498b0349c0..f486bd176c9f8 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -2534,7 +2534,7 @@ def AMDGPUMaxNumWorkGroups : InheritableAttr { def AMDGPUNamedBarrierWrapper : InheritableAttr { let Spellings = []; - let Args = [DeclArgument<Field, "WrappedField">]; + let Args = []; let Documentation = [InternalOnly]; let SemaHandler = 0; } diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 3ea0ed78fb3d3..a0aa64a489450 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -14298,12 +14298,15 @@ 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">; -def note_amdgcn_not_a_named_barrier_wrapper_derived_class : Note< - "%0 is not a named barrier wrapper because it inherits from %1; named barrier wrappers may not have any bases">; 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 note_amdgcn_named_barrier_wrapper_implies_final : Note< - "%0 is a named barrier wrapper around %1, which implies 'final'">; + +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">; diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp index b951c2f6552dd..6df6274d7e42f 100644 --- a/clang/lib/Sema/SemaAMDGPU.cpp +++ b/clang/lib/Sema/SemaAMDGPU.cpp @@ -1063,61 +1063,83 @@ void SemaAMDGPU::DiagnoseUnguardedBuiltinUsage(FunctionDecl *FD) { DiagnoseUnguardedBuiltins(SemaRef).IssueDiagnostics(FD->getBody()); } -void SemaAMDGPU::checkNamedBarrierWrapper(RecordDecl *R) { - if (R->isInvalidDecl()) - return; - - // Check if this record has any fields that interest us. - // We cannot codegen any kind of heterogenous struct where we named barrier - // fields mixed with other fields. - FieldDecl *NamedBarrField = nullptr; +static FieldDecl *getNamedBarrierField(const RecordDecl *R) { for (FieldDecl *FD : R->fields()) { QualType FDTy = FD->getType(); - if (FDTy->isAMDGPUNamedBarrierTypeOrWrapper()) { - NamedBarrField = FD; - break; - } + if (FDTy->isAMDGPUNamedBarrierTypeOrWrapper()) + return FD; } - if (!NamedBarrField) + return nullptr; +} + +void SemaAMDGPU::checkNamedBarrierWrapper(RecordDecl *R) { + if (R->isInvalidDecl()) return; - bool IsInvalid = false; - const auto OnError = [&]() { - if (!IsInvalid) { + 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(); - IsInvalid = true; + SemaRef.Diag( + R->getLocation(), + diag::note_amdgcn_not_a_named_barrier_wrapper_too_many_fields) + << R->getDeclName(); + return; } - }; - if (R->getNumFields() > 1) { - OnError(); - SemaRef.Diag(R->getLocation(), - diag::note_amdgcn_not_a_named_barrier_wrapper_too_many_fields) - << R->getName(); + IsWrapper = true; + DiagWrapperNote = [this, R, NamedBarrField]() { + SemaRef.Diag(NamedBarrField->getLocation(), + diag::note_amdgcn_named_barrier_reason_field) + << R->getDeclName() << NamedBarrField->getDeclName(); + }; } - if (const auto *CxxR = dyn_cast<CXXRecordDecl>(R)) { - if (CxxR->getNumBases() != 0) { - OnError(); - for (CXXBaseSpecifier CxxBase : CxxR->bases()) { - SemaRef.Diag( - CxxBase.getBaseTypeLoc(), - diag::note_amdgcn_not_a_named_barrier_wrapper_derived_class) - << R->getName() << CxxBase.getType(); - } + // 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 (IsInvalid) + if (!IsWrapper) return; + // Set the attribute even if the wrapper may be found to be invalid later. ASTContext &Context = getASTContext(); - SourceRange SR = NamedBarrField->getSourceRange(); - R->addAttr(FinalAttr::CreateImplicit(Context, SR)); - R->addAttr(AMDGPUNamedBarrierWrapperAttr::CreateImplicit(Context, - NamedBarrField, SR)); + 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/SemaDeclCXX.cpp b/clang/lib/Sema/SemaDeclCXX.cpp index 411e247523b73..685412d681680 100644 --- a/clang/lib/Sema/SemaDeclCXX.cpp +++ b/clang/lib/Sema/SemaDeclCXX.cpp @@ -2854,19 +2854,6 @@ CXXBaseSpecifier *Sema::CheckBaseSpecifier(CXXRecordDecl *Class, << BaseDecl->getDeclName() << FA->isSpelledAsSealed(); Diag(BaseDecl->getLocation(), diag::note_entity_declared_at) << BaseDecl->getDeclName() << FA->getRange(); - - // AMDGPU Named Barrier Wrappers are final. If the FinalAttr may have come - // from there, diagnose it. - if (FA->isImplicit()) { - if (auto *NBW = BaseDecl->getAttr<AMDGPUNamedBarrierWrapperAttr>()) { - FieldDecl *WrappedField = NBW->getWrappedField(); - Diag(BaseDecl->getLocation(), - diag::note_amdgcn_named_barrier_wrapper_implies_final) - << BaseDecl->getDeclName() << WrappedField->getDeclName(); - Diag(WrappedField->getLocation(), diag::note_entity_declared_at) - << WrappedField->getDeclName(); - } - } return nullptr; } diff --git a/clang/test/SemaCXX/amdgpu-barrier.cpp b/clang/test/SemaCXX/amdgpu-barrier.cpp index f9e5313bcfa26..80816fba6105e 100644 --- a/clang/test/SemaCXX/amdgpu-barrier.cpp +++ b/clang/test/SemaCXX/amdgpu-barrier.cpp @@ -15,10 +15,8 @@ void foo() { using SugaredArray = __amdgpu_named_workgroup_barrier_t[2]; -// expected-note@+2 {{'TestSimple' is a named barrier wrapper around 'x', which implies 'final'}} -// expected-note@+1 {{'TestSimple' declared here}} struct TestSimple { - __amdgpu_named_workgroup_barrier_t x; // expected-note {{'x' declared here}} + __amdgpu_named_workgroup_barrier_t x; }; struct TestArray{ @@ -35,27 +33,56 @@ 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}} +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 cannot have base class +// Wrappers must have standard layout struct WrapperBase { +__amdgpu_named_workgroup_barrier_t x; }; -struct WrapperWithBase : public WrapperBase { // expected-note {{WrapperWithBase is not a named barrier wrapper because it inherits from 'WrapperBase'; named barrier wrappers may not have any bases}} - __amdgpu_named_workgroup_barrier_t x; // expected-error {{fields of type '__amdgpu_named_workgroup_barrier_t' are only allowed in named barrier wrappers}} +struct WrapperWithBase : public WrapperBase { }; -// Wrappers imply final -struct NoWrapperInheritance : public TestSimple { }; // expected-error {{base 'TestSimple' is marked 'final'}} +// 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 final { // expected-note {{WrapperOfWrapperWithTooManyFields is not a named barrier wrapper because it has more than one field}} +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 a0baaaa3bf899..e8cd52d77d4a5 100644 --- a/clang/test/SemaHIP/amdgpu-barrier.hip +++ b/clang/test/SemaHIP/amdgpu-barrier.hip @@ -18,10 +18,8 @@ __device__ void foo() { using SugaredArray = __amdgpu_named_workgroup_barrier_t[2]; -// expected-note@+2 {{'TestSimple' is a named barrier wrapper around 'x', which implies 'final'}} -// expected-note@+1 {{'TestSimple' declared here}} struct TestSimple { - __amdgpu_named_workgroup_barrier_t x; // expected-note {{'x' declared here}} + __amdgpu_named_workgroup_barrier_t x; }; struct TestArray{ @@ -38,27 +36,57 @@ 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}} +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 cannot have base class +// Wrappers must have standard layout struct WrapperBase { +__amdgpu_named_workgroup_barrier_t x; }; -struct WrapperWithBase : public WrapperBase { // expected-note {{WrapperWithBase is not a named barrier wrapper because it inherits from 'WrapperBase'; named barrier wrappers may not have any bases}} - __amdgpu_named_workgroup_barrier_t x; // expected-error {{fields of type '__amdgpu_named_workgroup_barrier_t' are only allowed in named barrier wrappers}} +struct WrapperWithBase : public WrapperBase { }; -// Wrappers imply final -struct NoWrapperInheritance : public TestSimple { }; // expected-error {{base 'TestSimple' is marked 'final'}} +// 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 final { // expected-note {{WrapperOfWrapperWithTooManyFields is not a named barrier wrapper because it has more than one field}} +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 2f96a6aba3a3e..0898b1642f96f 100644 --- a/clang/test/SemaOpenCL/amdgpu-barrier.cl +++ b/clang/test/SemaOpenCL/amdgpu-barrier.cl @@ -22,7 +22,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}} + 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; }; >From 64b76dca15efbcad71d96cfa2a11f6368b67714f Mon Sep 17 00:00:00 2001 From: pvanhout <[email protected]> Date: Thu, 9 Jul 2026 11:00:32 +0200 Subject: [PATCH 07/10] Revert earlier change --- clang/include/clang/Basic/DiagnosticSemaKinds.td | 2 +- clang/lib/Sema/SemaDecl.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index a0aa64a489450..effa7516f9a5a 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -11885,7 +11885,7 @@ def note_within_param_of_type : Note< "within parameter %0 of type %1 declared here">; def note_illegal_field_declared_here : Note< "field of illegal %select{type|pointer type}0 %1 declared here">; -def err_invalid_type_for_struct_or_union_field : Error< +def err_opencl_type_struct_or_union_field : Error< "the %0 type cannot be used to declare a structure or union field">; def err_event_t_addr_space_qual : Error< "the event_t type can only be used with __private address space qualifier">; diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 635dd2d30e578..f71505800ac3f 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -19476,7 +19476,7 @@ FieldDecl *Sema::CheckFieldDecl(DeclarationName Name, QualType T, // used as structure or union field: image, sampler, event or block types. if (T->isEventT() || T->isImageType() || T->isSamplerT() || T->isBlockPointerType()) { - Diag(Loc, diag::err_invalid_type_for_struct_or_union_field) << T; + Diag(Loc, diag::err_opencl_type_struct_or_union_field) << T; Record->setInvalidDecl(); InvalidDecl = true; } >From d387f72a80d647dbb6bdcc02401ac810a288217a Mon Sep 17 00:00:00 2001 From: pvanhout <[email protected]> Date: Mon, 13 Jul 2026 16:56:40 +0200 Subject: [PATCH 08/10] Factor out check for AMDGPU types so it can be reused. --- clang/include/clang/Basic/TargetInfo.h | 7 +++++++ clang/lib/AST/ASTContext.cpp | 8 +------- clang/lib/Basic/TargetInfo.cpp | 2 ++ clang/lib/Basic/Targets/AMDGPU.cpp | 1 + clang/lib/Basic/Targets/SPIR.h | 1 + clang/lib/Sema/Sema.cpp | 9 ++------- clang/lib/Sema/SemaAMDGPU.cpp | 7 ++++++- 7 files changed, 20 insertions(+), 15 deletions(-) diff --git a/clang/include/clang/Basic/TargetInfo.h b/clang/include/clang/Basic/TargetInfo.h index 3aba4d261a651..1db2fd5698e97 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 ANDGPU 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/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/Basic/TargetInfo.cpp b/clang/lib/Basic/TargetInfo.cpp index e14e5f7489168..5b30fc2623ecf 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; @@ -168,6 +169,7 @@ TargetInfo::TargetInfo(const llvm::Triple &T) : Triple(T) { HasBuiltinZOSVaList = false; HasAArch64ACLETypes = false; HasRISCVVTypes = false; + HasAMDGPUTypes = (Triple.getVendor() == llvm::Triple::AMD); AllowAMDGPUUnsafeFPAtomics = false; HasUnalignedAccess = false; ARMCDECoprocMask = 0; diff --git a/clang/lib/Basic/Targets/AMDGPU.cpp b/clang/lib/Basic/Targets/AMDGPU.cpp index f9855161c0837..1e8a0f7f79ba5 100644 --- a/clang/lib/Basic/Targets/AMDGPU.cpp +++ b/clang/lib/Basic/Targets/AMDGPU.cpp @@ -201,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/SPIR.h b/clang/lib/Basic/Targets/SPIR.h index 0944dccb04e4a..1df6eb3c77975 100644 --- a/clang/lib/Basic/Targets/SPIR.h +++ b/clang/lib/Basic/Targets/SPIR.h @@ -312,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 = true; } llvm::SmallVector<Builtin::InfosShard> getTargetBuiltins() const override; 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 6df6274d7e42f..309afdb64b110 100644 --- a/clang/lib/Sema/SemaAMDGPU.cpp +++ b/clang/lib/Sema/SemaAMDGPU.cpp @@ -1074,9 +1074,15 @@ static FieldDecl *getNamedBarrierField(const RecordDecl *R) { } 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; @@ -1128,7 +1134,6 @@ void SemaAMDGPU::checkNamedBarrierWrapper(RecordDecl *R) { return; // Set the attribute even if the wrapper may be found to be invalid later. - ASTContext &Context = getASTContext(); R->addAttr( AMDGPUNamedBarrierWrapperAttr::CreateImplicit(Context, SourceRange())); >From bb668ba0810c4d3c172f1c540aa4540e4a433520 Mon Sep 17 00:00:00 2001 From: pvanhout <[email protected]> Date: Mon, 13 Jul 2026 16:59:26 +0200 Subject: [PATCH 09/10] Add RUN lines to check Sema check works if AMDGCN is aux target --- clang/test/SemaCXX/amdgpu-barrier.cpp | 1 + clang/test/SemaOpenCL/amdgpu-barrier.cl | 1 + 2 files changed, 2 insertions(+) diff --git a/clang/test/SemaCXX/amdgpu-barrier.cpp b/clang/test/SemaCXX/amdgpu-barrier.cpp index 80816fba6105e..d6869764ccfe3 100644 --- a/clang/test/SemaCXX/amdgpu-barrier.cpp +++ b/clang/test/SemaCXX/amdgpu-barrier.cpp @@ -1,5 +1,6 @@ // REQUIRES: amdgpu-registered-target // RUN: %clang_cc1 -fsyntax-only -verify -std=gnu++11 -triple amdgcn -Wno-unused-value %s +// RUN: %clang_cc1 -fcuda-is-device -fsyntax-only -verify -std=gnu++11 -triple x86_64 -aux-triple amdgcn -Wno-unused-value %s void foo() { int n = 100; diff --git a/clang/test/SemaOpenCL/amdgpu-barrier.cl b/clang/test/SemaOpenCL/amdgpu-barrier.cl index 0898b1642f96f..aabfc5e28fec2 100644 --- a/clang/test/SemaOpenCL/amdgpu-barrier.cl +++ b/clang/test/SemaOpenCL/amdgpu-barrier.cl @@ -1,6 +1,7 @@ // REQUIRES: amdgpu-registered-target // RUN: %clang_cc1 -verify -cl-std=CL1.2 -triple amdgcn-amd-amdhsa -Wno-unused-value %s // RUN: %clang_cc1 -verify -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -Wno-unused-value %s +// RUN: %clang_cc1 -fcuda-is-device -cl-std=CL2.0 -verify -triple x86_64 -aux-triple amdgcn -Wno-unused-value %s void foo() { typedef __amdgpu_named_workgroup_barrier_t SugaredArray[2]; >From fbb063805463671de716f81e7ebecbb478b6e717 Mon Sep 17 00:00:00 2001 From: pvanhout <[email protected]> Date: Tue, 14 Jul 2026 10:24:05 +0200 Subject: [PATCH 10/10] Address Comments --- clang/include/clang/Basic/TargetInfo.h | 2 +- clang/lib/Basic/TargetInfo.cpp | 2 +- clang/lib/Basic/Targets/SPIR.h | 2 +- clang/test/CodeGenHIP/amdgpu-barrier-type.hip | 37 +++++++------------ clang/test/SemaCXX/amdgpu-barrier.cpp | 1 - clang/test/SemaOpenCL/amdgpu-barrier.cl | 1 - 6 files changed, 16 insertions(+), 29 deletions(-) diff --git a/clang/include/clang/Basic/TargetInfo.h b/clang/include/clang/Basic/TargetInfo.h index 1db2fd5698e97..871666af1da77 100644 --- a/clang/include/clang/Basic/TargetInfo.h +++ b/clang/include/clang/Basic/TargetInfo.h @@ -1076,7 +1076,7 @@ class TargetInfo : public TransferrableTargetInfo, /// available on this target. bool hasAArch64ACLETypes() const { return HasAArch64ACLETypes; } - /// Returns whether or not the ANDGPU built-in types are + /// Returns whether or not the AMDGPU built-in types are /// available on this target. bool hasAMDGPUTypes() const { return HasAMDGPUTypes; } diff --git a/clang/lib/Basic/TargetInfo.cpp b/clang/lib/Basic/TargetInfo.cpp index 5b30fc2623ecf..e0b745f8f23ac 100644 --- a/clang/lib/Basic/TargetInfo.cpp +++ b/clang/lib/Basic/TargetInfo.cpp @@ -169,7 +169,7 @@ TargetInfo::TargetInfo(const llvm::Triple &T) : Triple(T) { HasBuiltinZOSVaList = false; HasAArch64ACLETypes = false; HasRISCVVTypes = false; - HasAMDGPUTypes = (Triple.getVendor() == llvm::Triple::AMD); + HasAMDGPUTypes = false; AllowAMDGPUUnsafeFPAtomics = false; HasUnalignedAccess = false; ARMCDECoprocMask = 0; diff --git a/clang/lib/Basic/Targets/SPIR.h b/clang/lib/Basic/Targets/SPIR.h index 1df6eb3c77975..37baf7c99a90c 100644 --- a/clang/lib/Basic/Targets/SPIR.h +++ b/clang/lib/Basic/Targets/SPIR.h @@ -312,7 +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 = true; + HasAMDGPUTypes = (Triple.getVendor() == llvm::Triple::AMD); } llvm::SmallVector<Builtin::InfosShard> getTargetBuiltins() const override; diff --git a/clang/test/CodeGenHIP/amdgpu-barrier-type.hip b/clang/test/CodeGenHIP/amdgpu-barrier-type.hip index 3664821b80541..df9e3631c0d1f 100644 --- a/clang/test/CodeGenHIP/amdgpu-barrier-type.hip +++ b/clang/test/CodeGenHIP/amdgpu-barrier-type.hip @@ -1,27 +1,25 @@ -// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --check-globals +// 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__ __amdgpu_named_workgroup_barrier_t bar_arr[2]; //. // CHECK: @bar = addrspace(3) global target("amdgcn.named.barrier", 0) undef, align 4 -// CHECK: @arr = addrspace(3) global [2 x target("amdgcn.named.barrier", 0)] undef, align 4 -// CHECK: @wrapper_str = addrspace(3) global %struct.WrapperStruct undef, align 4 -// CHECK: @wrapperwrapper_str = addrspace(3) global %struct.WrapperWrapperStruct undef, align 4 -// CHECK: @__hip_cuid_ = addrspace(1) global i8 0 -// CHECK: @llvm.compiler.used = appending addrspace(1) global [1 x ptr] [ptr addrspacecast (ptr addrspace(1) @__hip_cuid_ to ptr)], section "llvm.metadata" +// 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; -} wrapper_str; +} bar_wrapper_str; __shared__ struct WrapperWrapperStruct { WrapperStruct x; -} wrapperwrapper_str; +} bar_wrapperwrapper_str; __attribute__((device)) __amdgpu_named_workgroup_barrier_t *getBar(); __attribute__((device)) void useBar(__amdgpu_named_workgroup_barrier_t *); @@ -35,9 +33,9 @@ __attribute__((device)) void useBar(__amdgpu_named_workgroup_barrier_t *); // 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(3) @bar to ptr)) #[[ATTR2]] -// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef addrspacecast (ptr addrspace(3) @wrapper_str to ptr)) #[[ATTR2]] -// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef addrspacecast (ptr addrspace(3) @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) @arr to ptr), i64 16)) #[[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]] @@ -46,18 +44,9 @@ __attribute__((device)) void useBar(__amdgpu_named_workgroup_barrier_t *); __attribute__((device)) __amdgpu_named_workgroup_barrier_t *testSem(__amdgpu_named_workgroup_barrier_t *p) { useBar(p); useBar(&bar); - useBar(&wrapper_str.x); - useBar(&wrapperwrapper_str.x.x); - useBar(&arr[1]); + useBar(&bar_wrapper_str.x); + useBar(&bar_wrapperwrapper_str.x.x); + useBar(&bar_arr[1]); useBar(getBar()); return getBar(); } -//. -// CHECK: attributes #[[ATTR0]] = { convergent mustprogress noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="gfx1250" "uniform-work-group-size" } -// CHECK: attributes #[[ATTR1:[0-9]+]] = { convergent nounwind "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="gfx1250" "uniform-work-group-size" } -// CHECK: attributes #[[ATTR2]] = { convergent nounwind "uniform-work-group-size" } -//. -// CHECK: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 600} -// CHECK: [[META1:![0-9]+]] = !{i32 1, !"amdgpu_printf_kind", !"hostcall"} -// CHECK: [[META2:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"} -//. diff --git a/clang/test/SemaCXX/amdgpu-barrier.cpp b/clang/test/SemaCXX/amdgpu-barrier.cpp index d6869764ccfe3..80816fba6105e 100644 --- a/clang/test/SemaCXX/amdgpu-barrier.cpp +++ b/clang/test/SemaCXX/amdgpu-barrier.cpp @@ -1,6 +1,5 @@ // REQUIRES: amdgpu-registered-target // RUN: %clang_cc1 -fsyntax-only -verify -std=gnu++11 -triple amdgcn -Wno-unused-value %s -// RUN: %clang_cc1 -fcuda-is-device -fsyntax-only -verify -std=gnu++11 -triple x86_64 -aux-triple amdgcn -Wno-unused-value %s void foo() { int n = 100; diff --git a/clang/test/SemaOpenCL/amdgpu-barrier.cl b/clang/test/SemaOpenCL/amdgpu-barrier.cl index aabfc5e28fec2..0898b1642f96f 100644 --- a/clang/test/SemaOpenCL/amdgpu-barrier.cl +++ b/clang/test/SemaOpenCL/amdgpu-barrier.cl @@ -1,7 +1,6 @@ // REQUIRES: amdgpu-registered-target // RUN: %clang_cc1 -verify -cl-std=CL1.2 -triple amdgcn-amd-amdhsa -Wno-unused-value %s // RUN: %clang_cc1 -verify -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -Wno-unused-value %s -// RUN: %clang_cc1 -fcuda-is-device -cl-std=CL2.0 -verify -triple x86_64 -aux-triple amdgcn -Wno-unused-value %s void foo() { typedef __amdgpu_named_workgroup_barrier_t SugaredArray[2]; _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
