https://github.com/Pierre-vh updated https://github.com/llvm/llvm-project/pull/207687
>From 0efe9d7a5777dced26a75ebebc012499a88a10bf Mon Sep 17 00:00:00 2001 From: pvanhout <[email protected]> Date: Mon, 6 Jul 2026 11:33:34 +0200 Subject: [PATCH 1/2] [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 3a801e2857b13..48612598735b9 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 86b765fdf1fab..437f0accbb8b5 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -11855,7 +11855,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 19b85d0e0af69..51417d553f29f 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 9a25384347073..5e8871ace9a60 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 fc8e0e6071711..2974ddf551398 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 8b21b86bac264..6c4721c22e9ee 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 8af9cb91c47ab..8f3638578b264 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 389cc075a3a0b..5ffd427c77f91 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 70b529e8ca854..bb8fa1a1b3dec 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 88192a756cc4f..ab498a082302b 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 b0daffbea4576..e426b9113f38a 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 e155fdd752d7f..7fd54c6c34cb7 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -6171,6 +6171,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 9948a633d7e98..b77305d575e67 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -19435,7 +19435,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; } @@ -19448,6 +19448,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 cf4ab5958bfd190d774d6eed0043c9b5ce6782a2 Mon Sep 17 00:00:00 2001 From: pvanhout <[email protected]> Date: Mon, 6 Jul 2026 13:56:08 +0200 Subject: [PATCH 2/2] 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 51417d553f29f..ab7653d457739 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 5e8871ace9a60..d4249fb065375 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 2974ddf551398..1faae0b213bf6 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 6c4721c22e9ee..7c1aa7c75f726 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 8f3638578b264..0d452b7f33bd3 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 5ffd427c77f91..f9153986b5fe1 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 bb8fa1a1b3dec..7b08638d7232c 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 ab498a082302b..b60963abbceaf 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 e426b9113f38a..35c8227de8d27 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 7fd54c6c34cb7..56c39bef1b5ef 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -6171,11 +6171,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; _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
