[clang] [clang][AMDGPU] Clean-up handling of named barrier type (PR #207687)
Pierre van Houtryve via cfe-commits
cfe-commits at lists.llvm.org
Tue Jul 14 01:24:28 PDT 2026
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 <pierre.vanhoutryve at amd.com>
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 <pierre.vanhoutryve at amd.com>
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 <pierre.vanhoutryve at amd.com>
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 <pierre.vanhoutryve at amd.com>
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 <pierre.vanhoutryve at amd.com>
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 at +2 {{'TestSimple' is a named barrier wrapper around 'x', which implies 'final'}}
+// expected-note at +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 at +2 {{'TestSimple' is a named barrier wrapper around 'x', which implies 'final'}}
+// expected-note at +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 <pierre.vanhoutryve at amd.com>
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 at +2 {{'TestSimple' is a named barrier wrapper around 'x', which implies 'final'}}
-// expected-note at +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 at +2 {{named barrier wrapper 'WrapperWithBaseNoStandardLayout' must have a C++11 standard layou}}
+// expected-note at +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 at +2 {{named barrier wrapper 'TemplatedWrapperImplNoStandardLayout' must have a C++11 standard layou}}
+// expected-note at +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 at +2 {{'TestSimple' is a named barrier wrapper around 'x', which implies 'final'}}
-// expected-note at +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 at +2 {{named barrier wrapper 'WrapperWithBaseNoStandardLayout' must have a C++11 standard layou}}
+// expected-note at +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 at +2 {{named barrier wrapper 'TemplatedWrapperImplNoStandardLayout' must have a C++11 standard layou}}
+// expected-note at +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 <pierre.vanhoutryve at amd.com>
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 <pierre.vanhoutryve at amd.com>
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 <pierre.vanhoutryve at amd.com>
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 <pierre.vanhoutryve at amd.com>
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];
More information about the cfe-commits
mailing list