Skip to content

Commit 4ac0e7e

Browse files
cmc-repcmc-rep
authored
[AMDGPU] Add a type for the named barrier (#113614)
1 parent 88cc7ac commit 4ac0e7e

File tree

12 files changed

+168
-4
lines changed

12 files changed

+168
-4
lines changed

‎clang/include/clang/Basic/AMDGPUTypes.def

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -15,7 +15,15 @@
1515
AMDGPU_TYPE(Name, Id, SingletonId, Width, Align)
1616
#endif
1717

18+
#ifndef AMDGPU_NAMED_BARRIER_TYPE
19+
#define AMDGPU_NAMED_BARRIER_TYPE(Name, Id, SingletonId, Width, Align, Scope) \
20+
AMDGPU_TYPE(Name, Id, SingletonId, Width, Align)
21+
#endif
22+
1823
AMDGPU_OPAQUE_PTR_TYPE("__amdgpu_buffer_rsrc_t", AMDGPUBufferRsrc, AMDGPUBufferRsrcTy, 128, 128, 8)
1924

25+
AMDGPU_NAMED_BARRIER_TYPE("__amdgpu_named_workgroup_barrier_t", AMDGPUNamedWorkgroupBarrier, AMDGPUNamedWorkgroupBarrierTy, 128, 32, 0)
26+
2027
#undef AMDGPU_TYPE
2128
#undef AMDGPU_OPAQUE_PTR_TYPE
29+
#undef AMDGPU_NAMED_BARRIER_TYPE

‎clang/lib/CodeGen/CGDebugInfo.cpp

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -916,6 +916,13 @@ llvm::DIType *CGDebugInfo::CreateType(const BuiltinType *BT) {
916916
TheCU, TheCU->getFile(), 0); \
917917
return SingletonId; \
918918
}
919+
#define AMDGPU_NAMED_BARRIER_TYPE(Name, Id, SingletonId, Width, Align, Scope) \
920+
case BuiltinType::Id: { \
921+
if (!SingletonId) \
922+
SingletonId = \
923+
DBuilder.createBasicType(Name, Width, llvm::dwarf::DW_ATE_unsigned); \
924+
return SingletonId; \
925+
}
919926
#include "clang/Basic/AMDGPUTypes.def"
920927
case BuiltinType::UChar:
921928
case BuiltinType::Char_U:

‎clang/lib/CodeGen/CodeGenTypes.cpp

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -564,6 +564,10 @@ llvm::Type *CodeGenTypes::ConvertType(QualType T) {
564564
#define AMDGPU_OPAQUE_PTR_TYPE(Name, Id, SingletonId, Width, Align, AS) \
565565
case BuiltinType::Id: \
566566
return llvm::PointerType::get(getLLVMContext(), AS);
567+
#define AMDGPU_NAMED_BARRIER_TYPE(Name, Id, SingletonId, Width, Align, Scope) \
568+
case BuiltinType::Id: \
569+
return llvm::TargetExtType::get(getLLVMContext(), "amdgcn.named.barrier", \
570+
{}, {Scope});
567571
#include "clang/Basic/AMDGPUTypes.def"
568572
#define HLSL_INTANGIBLE_TYPE(Name, Id, SingletonId) case BuiltinType::Id:
569573
#include "clang/Basic/HLSLIntangibleTypes.def"

‎clang/test/AST/ast-dump-amdgpu-types.c

Lines changed: 9 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1,10 +1,15 @@
11
// REQUIRES: amdgpu-registered-target
22
// Test without serialization:
3-
// RUN: %clang_cc1 -triple amdgcn -ast-dump -ast-dump-filter __amdgpu_buffer_rsrc_t %s | FileCheck %s
3+
// RUN: %clang_cc1 -triple amdgcn -ast-dump -ast-dump-filter __amdgpu_buffer_rsrc_t %s | FileCheck %s -check-prefix=BUFFER-RSRC
4+
// RUN: %clang_cc1 -triple amdgcn -ast-dump -ast-dump-filter __amdgpu_named_workgroup_barrier %s | FileCheck %s -check-prefix=WORKGROUP-BARRIER
45
//
56
// Test with serialization:
67
// RUN: %clang_cc1 -triple amdgcn -emit-pch -o %t %s
7-
// RUN: %clang_cc1 -x c -triple amdgcn -include-pch %t -ast-dump-all -ast-dump-filter __amdgpu_buffer_rsrc_t /dev/null | sed -e "s/ <undeserialized declarations>//" -e "s/ imported//" | FileCheck %s
8+
// RUN: %clang_cc1 -x c -triple amdgcn -include-pch %t -ast-dump-all -ast-dump-filter __amdgpu_buffer_rsrc_t /dev/null | sed -e "s/ <undeserialized declarations>//" -e "s/ imported//" | FileCheck %s -check-prefix=BUFFER-RSRC
9+
// RUN: %clang_cc1 -x c -triple amdgcn -include-pch %t -ast-dump-all -ast-dump-filter __amdgpu_named_workgroup_barrier /dev/null | sed -e "s/ <undeserialized declarations>//" -e "s/ imported//" | FileCheck %s -check-prefix=WORKGROUP-BARRIER
810

9-
// CHECK: TypedefDecl {{.*}} implicit __amdgpu_buffer_rsrc_t
10-
// CHECK-NEXT: -BuiltinType {{.*}} '__amdgpu_buffer_rsrc_t'
11+
// BUFFER-RSRC: TypedefDecl {{.*}} implicit __amdgpu_buffer_rsrc_t
12+
// BUFFER-RSRC-NEXT: -BuiltinType {{.*}} '__amdgpu_buffer_rsrc_t'
13+
14+
// WORKGROUP-BARRIER: TypedefDecl {{.*}} implicit __amdgpu_named_workgroup_barrier_t
15+
// WORKGROUP-BARRIER-NEXT: -BuiltinType {{.*}} '__amdgpu_named_workgroup_barrier_t'

‎clang/test/CodeGen/amdgpu-barrier-type-debug-info.c

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,8 @@
1+
// REQUIRES: amdgpu-registered-target
2+
// RUN: %clang_cc1 -triple amdgcn -emit-llvm -o - %s -debug-info-kind=limited 2>&1 | FileCheck %s
3+
4+
// CHECK: name: "__amdgpu_named_workgroup_barrier_t",{{.*}}baseType: ![[BT:[0-9]+]]
5+
// CHECK: [[BT]] = !DIBasicType(name: "__amdgpu_named_workgroup_barrier_t", size: 128, encoding: DW_ATE_unsigned)
6+
void test_locals(void) {
7+
__amdgpu_named_workgroup_barrier_t k0;
8+
}

‎clang/test/CodeGenCXX/amdgpu-barrier-typeinfo.cpp

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,10 @@
1+
// REQUIRES: amdgpu-registered-target
2+
// RUN: %clang_cc1 -triple amdgcn %s -emit-llvm -o - | FileCheck %s
3+
4+
namespace std { class type_info; };
5+
6+
auto &b0 = typeid(__amdgpu_named_workgroup_barrier_t);
7+
8+
// CHECK-DAG: @_ZTSu34__amdgpu_named_workgroup_barrier_t = {{.*}} c"u34__amdgpu_named_workgroup_barrier_t\00"
9+
// CHECK-DAG: @_ZTIu34__amdgpu_named_workgroup_barrier_t = {{.*}} @_ZTVN10__cxxabiv123__fundamental_type_infoE, {{.*}} @_ZTSu34__amdgpu_named_workgroup_barrier_t
10+

‎clang/test/CodeGenHIP/amdgpu-barrier-type.hip

Lines changed: 42 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,42 @@
1+
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature
2+
// REQUIRES: amdgpu-registered-target
3+
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu verde -emit-llvm -o - %s | FileCheck %s
4+
5+
#define __shared__ __attribute__((shared))
6+
7+
__shared__ __amdgpu_named_workgroup_barrier_t bar;
8+
__shared__ __amdgpu_named_workgroup_barrier_t arr[2];
9+
__shared__ struct {
10+
__amdgpu_named_workgroup_barrier_t x;
11+
__amdgpu_named_workgroup_barrier_t y;
12+
} str;
13+
14+
__amdgpu_named_workgroup_barrier_t *getBar();
15+
void useBar(__amdgpu_named_workgroup_barrier_t *);
16+
17+
// CHECK-LABEL: define {{[^@]+}}@_Z7testSemPu34__amdgpu_named_workgroup_barrier_t
18+
// CHECK-SAME: (ptr noundef [[P:%.*]]) #[[ATTR0:[0-9]+]] {
19+
// CHECK-NEXT: entry:
20+
// CHECK-NEXT: [[RETVAL:%.*]] = alloca ptr, align 8, addrspace(5)
21+
// CHECK-NEXT: [[P_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
22+
// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
23+
// CHECK-NEXT: [[P_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[P_ADDR]] to ptr
24+
// CHECK-NEXT: store ptr [[P]], ptr [[P_ADDR_ASCAST]], align 8
25+
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[P_ADDR_ASCAST]], align 8
26+
// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef [[TMP0]]) #[[ATTR2:[0-9]+]]
27+
// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef addrspacecast (ptr addrspace(1) @bar to ptr)) #[[ATTR2]]
28+
// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef getelementptr inbounds ([2 x target("amdgcn.named.barrier", 0)], ptr addrspacecast (ptr addrspace(1) @arr to ptr), i64 0, i64 1)) #[[ATTR2]]
29+
// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef getelementptr inbounds nuw ([[STRUCT_ANON:%.*]], ptr addrspacecast (ptr addrspace(1) @str to ptr), i32 0, i32 1)) #[[ATTR2]]
30+
// CHECK-NEXT: [[CALL:%.*]] = call noundef ptr @_Z6getBarv() #[[ATTR2]]
31+
// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef [[CALL]]) #[[ATTR2]]
32+
// CHECK-NEXT: [[CALL1:%.*]] = call noundef ptr @_Z6getBarv() #[[ATTR2]]
33+
// CHECK-NEXT: ret ptr [[CALL1]]
34+
//
35+
__amdgpu_named_workgroup_barrier_t *testSem(__amdgpu_named_workgroup_barrier_t *p) {
36+
useBar(p);
37+
useBar(&bar);
38+
useBar(&arr[1]);
39+
useBar(&str.y);
40+
useBar(getBar());
41+
return getBar();
42+
}

‎clang/test/SemaCXX/amdgpu-barrier.cpp

Lines changed: 17 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,17 @@
1+
// REQUIRES: amdgpu-registered-target
2+
// RUN: %clang_cc1 -fsyntax-only -verify -std=gnu++11 -triple amdgcn -Wno-unused-value %s
3+
4+
void foo() {
5+
int n = 100;
6+
__amdgpu_named_workgroup_barrier_t v = 0; // expected-error {{cannot initialize a variable of type '__amdgpu_named_workgroup_barrier_t' with an rvalue of type 'int'}}
7+
static_cast<__amdgpu_named_workgroup_barrier_t>(n); // expected-error {{static_cast from 'int' to '__amdgpu_named_workgroup_barrier_t' is not allowed}}
8+
dynamic_cast<__amdgpu_named_workgroup_barrier_t>(n); // expected-error {{invalid target type '__amdgpu_named_workgroup_barrier_t' for dynamic_cast; target type must be a reference or pointer type to a defined class}}
9+
reinterpret_cast<__amdgpu_named_workgroup_barrier_t>(n); // expected-error {{reinterpret_cast from 'int' to '__amdgpu_named_workgroup_barrier_t' is not allowed}}
10+
int c(v); // expected-error {{cannot initialize a variable of type 'int' with an lvalue of type '__amdgpu_named_workgroup_barrier_t'}}
11+
__amdgpu_named_workgroup_barrier_t k;
12+
int *ip = (int *)k; // expected-error {{cannot cast from type '__amdgpu_named_workgroup_barrier_t' to pointer type 'int *'}}
13+
void *vp = (void *)k; // expected-error {{cannot cast from type '__amdgpu_named_workgroup_barrier_t' to pointer type 'void *'}}
14+
}
15+
16+
static_assert(sizeof(__amdgpu_named_workgroup_barrier_t) == 16, "wrong size");
17+
static_assert(alignof(__amdgpu_named_workgroup_barrier_t) == 4, "wrong alignment");

‎clang/test/SemaHIP/amdgpu-barrier.hip

Lines changed: 20 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,20 @@
1+
// REQUIRES: amdgpu-registered-target
2+
// RUN: %clang_cc1 -fsyntax-only -verify -triple amdgcn -Wno-unused-value %s
3+
// RUN: %clang_cc1 -fsyntax-only -verify -triple x86_64 -aux-triple amdgcn -Wno-unused-value %s
4+
5+
#define __device__ __attribute__((device))
6+
7+
__device__ void foo() {
8+
int n = 100;
9+
__amdgpu_named_workgroup_barrier_t v = 0; // expected-error {{cannot initialize a variable of type '__amdgpu_named_workgroup_barrier_t' with an rvalue of type 'int'}}
10+
static_cast<__amdgpu_named_workgroup_barrier_t>(n); // expected-error {{static_cast from 'int' to '__amdgpu_named_workgroup_barrier_t' is not allowed}}
11+
dynamic_cast<__amdgpu_named_workgroup_barrier_t>(n); // expected-error {{invalid target type '__amdgpu_named_workgroup_barrier_t' for dynamic_cast; target type must be a reference or pointer type to a defined class}}
12+
reinterpret_cast<__amdgpu_named_workgroup_barrier_t>(n); // expected-error {{reinterpret_cast from 'int' to '__amdgpu_named_workgroup_barrier_t' is not allowed}}
13+
int c(v); // expected-error {{cannot initialize a variable of type 'int' with an lvalue of type '__amdgpu_named_workgroup_barrier_t'}}
14+
__amdgpu_named_workgroup_barrier_t k;
15+
int *ip = (int *)k; // expected-error {{cannot cast from type '__amdgpu_named_workgroup_barrier_t' to pointer type 'int *'}}
16+
void *vp = (void *)k; // expected-error {{cannot cast from type '__amdgpu_named_workgroup_barrier_t' to pointer type 'void *'}}
17+
}
18+
19+
static_assert(sizeof(__amdgpu_named_workgroup_barrier_t) == 16, "wrong size");
20+
static_assert(alignof(__amdgpu_named_workgroup_barrier_t) == 4, "wrong alignment");

‎clang/test/SemaOpenCL/amdgpu-barrier.cl

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,12 @@
1+
// REQUIRES: amdgpu-registered-target
2+
// RUN: %clang_cc1 -verify -cl-std=CL1.2 -triple amdgcn-amd-amdhsa -Wno-unused-value %s
3+
// RUN: %clang_cc1 -verify -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -Wno-unused-value %s
4+
5+
void foo() {
6+
int n = 100;
7+
__amdgpu_named_workgroup_barrier_t v = 0; // expected-error {{initializing '__private __amdgpu_named_workgroup_barrier_t' with an expression of incompatible type 'int'}}
8+
int c = v; // expected-error {{initializing '__private int' with an expression of incompatible type '__private __amdgpu_named_workgroup_barrier_t'}}
9+
__amdgpu_named_workgroup_barrier_t k;
10+
int *ip = (int *)k; // expected-error {{operand of type '__amdgpu_named_workgroup_barrier_t' where arithmetic or pointer type is required}}
11+
void *vp = (void *)k; // expected-error {{operand of type '__amdgpu_named_workgroup_barrier_t' where arithmetic or pointer type is required}}
12+
}

‎clang/test/SemaOpenMP/amdgpu-barrier.cpp

Lines changed: 17 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,17 @@
1+
// REQUIRES: amdgpu-registered-target
2+
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -Wno-unused-value %s
3+
4+
void foo() {
5+
#pragma omp target
6+
{
7+
int n = 100;
8+
__amdgpu_named_workgroup_barrier_t v = 0; // expected-error {{cannot initialize a variable of type '__amdgpu_named_workgroup_barrier_t' with an rvalue of type 'int'}}
9+
static_cast<__amdgpu_named_workgroup_barrier_t>(n); // expected-error {{static_cast from 'int' to '__amdgpu_named_workgroup_barrier_t' is not allowed}}
10+
dynamic_cast<__amdgpu_named_workgroup_barrier_t>(n); // expected-error {{invalid target type '__amdgpu_named_workgroup_barrier_t' for dynamic_cast; target type must be a reference or pointer type to a defined class}}
11+
reinterpret_cast<__amdgpu_named_workgroup_barrier_t>(n); // expected-error {{reinterpret_cast from 'int' to '__amdgpu_named_workgroup_barrier_t' is not allowed}}
12+
int c(v); // expected-error {{cannot initialize a variable of type 'int' with an lvalue of type '__amdgpu_named_workgroup_barrier_t'}}
13+
__amdgpu_named_workgroup_barrier_t k;
14+
int *ip = (int *)k; // expected-error {{cannot cast from type '__amdgpu_named_workgroup_barrier_t' to pointer type 'int *'}}
15+
void *vp = (void *)k; // expected-error {{cannot cast from type '__amdgpu_named_workgroup_barrier_t' to pointer type 'void *'}}
16+
}
17+
}

‎llvm/lib/IR/Type.cpp

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -834,6 +834,14 @@ Expected<TargetExtType *> TargetExtType::checkParams(TargetExtType *TTy) {
834834
"target extension type riscv.vector.tuple should have one "
835835
"type parameter and one integer parameter");
836836

837+
// Opaque types in the AMDGPU name space.
838+
if (TTy->Name == "amdgcn.named.barrier" &&
839+
(TTy->getNumTypeParameters() != 0 || TTy->getNumIntParameters() != 1)) {
840+
return createStringError("target extension type amdgcn.named.barrier "
841+
"should have no type parameters "
842+
"and one integer parameter");
843+
}
844+
837845
return TTy;
838846
}
839847

@@ -879,6 +887,12 @@ static TargetTypeInfo getTargetTypeInfo(const TargetExtType *Ty) {
879887
if (Name.starts_with("dx."))
880888
return TargetTypeInfo(PointerType::get(C, 0));
881889

890+
// Opaque types in the AMDGPU name space.
891+
if (Name == "amdgcn.named.barrier") {
892+
return TargetTypeInfo(FixedVectorType::get(Type::getInt32Ty(C), 4),
893+
TargetExtType::CanBeGlobal);
894+
}
895+
882896
return TargetTypeInfo(Type::getVoidTy(C));
883897
}
884898

0 commit comments

Comments
 (0)