llvmorg-github-actions[bot] wrote:

<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-clang-codegen

Author: Pierre van Houtryve (Pierre-vh)

<details>
<summary>Changes</summary>

- 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

---
Full diff: https://github.com/llvm/llvm-project/pull/207687.diff


22 Files Affected:

- (modified) clang/include/clang/AST/TypeBase.h (+3) 
- (modified) clang/include/clang/Basic/AddressSpaces.h (+3) 
- (modified) clang/include/clang/Basic/DiagnosticSemaKinds.td (+1-1) 
- (modified) clang/lib/AST/Type.cpp (+13-1) 
- (modified) clang/lib/AST/TypePrinter.cpp (+2) 
- (modified) clang/lib/Basic/TargetInfo.cpp (+1) 
- (modified) clang/lib/Basic/Targets/AArch64.h (+1) 
- (modified) clang/lib/Basic/Targets/AMDGPU.cpp (+2) 
- (modified) clang/lib/Basic/Targets/DirectX.h (+1) 
- (modified) clang/lib/Basic/Targets/NVPTX.h (+1) 
- (modified) clang/lib/Basic/Targets/SPIR.h (+2) 
- (modified) clang/lib/Basic/Targets/SystemZ.h (+2-1) 
- (modified) clang/lib/Basic/Targets/TCE.h (+1) 
- (modified) clang/lib/Basic/Targets/WebAssembly.h (+1) 
- (modified) clang/lib/Basic/Targets/X86.h (+1) 
- (modified) clang/lib/CodeGen/CodeGenModule.cpp (+6) 
- (modified) clang/lib/Sema/SemaDecl.cpp (+9-1) 
- (modified) clang/test/CodeGenHIP/amdgpu-barrier-type.hip (+23-14) 
- (modified) clang/test/SemaCXX/amdgpu-barrier.cpp (+4) 
- (modified) clang/test/SemaHIP/amdgpu-barrier.hip (+5) 
- (modified) clang/test/SemaOpenCL/amdgpu-barrier.cl (+5) 
- (modified) clang/test/SemaTemplate/address_space-dependent.cpp (+2-2) 


``````````diff
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;

``````````

</details>


https://github.com/llvm/llvm-project/pull/207687
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits

Reply via email to