[clang] [llvm] [AMDGPU] Add a type for the named barrier (PR #113614)
Gang Chen via cfe-commits
cfe-commits at lists.llvm.org
Thu Oct 24 16:12:12 PDT 2024
https://github.com/cmc-rep updated https://github.com/llvm/llvm-project/pull/113614
>From 166a4aec8a8ee813be0ee3045563cd45efd944c0 Mon Sep 17 00:00:00 2001
From: gangc <gangc at amd.com>
Date: Thu, 24 Oct 2024 11:18:22 -0700
Subject: [PATCH 1/3] [AMDGPU] Add a type for the named barrier
---
clang/include/clang/Basic/AMDGPUTypes.def | 8 ++++
.../include/clang/Serialization/ASTBitCodes.h | 2 +-
clang/lib/CodeGen/CGDebugInfo.cpp | 7 ++++
clang/lib/CodeGen/CodeGenTypes.cpp | 4 ++
clang/test/AST/ast-dump-amdgpu-types.c | 13 ++++--
.../CodeGen/amdgpu-barrier-type-debug-info.c | 8 ++++
.../CodeGenCXX/amdgpu-barrier-typeinfo.cpp | 10 +++++
clang/test/CodeGenHIP/amdgpu-barrier-type.hip | 42 +++++++++++++++++++
clang/test/SemaCXX/amdgpu-barrier.cpp | 17 ++++++++
clang/test/SemaHIP/amdgpu-barrier.hip | 20 +++++++++
clang/test/SemaOpenCL/amdgpu-barrier.cl | 12 ++++++
clang/test/SemaOpenMP/amdgpu-barrier.cpp | 17 ++++++++
llvm/lib/IR/Type.cpp | 8 ++++
13 files changed, 163 insertions(+), 5 deletions(-)
create mode 100644 clang/test/CodeGen/amdgpu-barrier-type-debug-info.c
create mode 100644 clang/test/CodeGenCXX/amdgpu-barrier-typeinfo.cpp
create mode 100644 clang/test/CodeGenHIP/amdgpu-barrier-type.hip
create mode 100644 clang/test/SemaCXX/amdgpu-barrier.cpp
create mode 100644 clang/test/SemaHIP/amdgpu-barrier.hip
create mode 100644 clang/test/SemaOpenCL/amdgpu-barrier.cl
create mode 100644 clang/test/SemaOpenMP/amdgpu-barrier.cpp
diff --git a/clang/include/clang/Basic/AMDGPUTypes.def b/clang/include/clang/Basic/AMDGPUTypes.def
index e47e544fdc82c1..6b98e311b4cf55 100644
--- a/clang/include/clang/Basic/AMDGPUTypes.def
+++ b/clang/include/clang/Basic/AMDGPUTypes.def
@@ -15,7 +15,15 @@
AMDGPU_TYPE(Name, Id, SingletonId, Width, Align)
#endif
+#ifndef AMDGPU_NAMED_BARRIER_TYPE
+#define AMDGPU_NAMED_BARRIER_TYPE(Name, Id, SingletonId, Width, Align, Scope) \
+ AMDGPU_TYPE(Name, Id, SingletonId, Width, Align)
+#endif
+
AMDGPU_OPAQUE_PTR_TYPE("__amdgpu_buffer_rsrc_t", AMDGPUBufferRsrc, AMDGPUBufferRsrcTy, 128, 128, 8)
+AMDGPU_NAMED_BARRIER_TYPE("__amdgpu_named_workgroup_barrier_t", AMDGPUNamedWorkgroupBarrier, AMDGPUNamedWorkgroupBarrierTy, 128, 32, 0)
+
#undef AMDGPU_TYPE
#undef AMDGPU_OPAQUE_PTR_TYPE
+#undef AMDGPU_NAMED_BARRIER_TYPE
\ No newline at end of file
diff --git a/clang/include/clang/Serialization/ASTBitCodes.h b/clang/include/clang/Serialization/ASTBitCodes.h
index 13173dc96e71ae..99232fd2135790 100644
--- a/clang/include/clang/Serialization/ASTBitCodes.h
+++ b/clang/include/clang/Serialization/ASTBitCodes.h
@@ -1149,7 +1149,7 @@ enum PredefinedTypeIDs {
///
/// Type IDs for non-predefined types will start at
/// NUM_PREDEF_TYPE_IDs.
-const unsigned NUM_PREDEF_TYPE_IDS = 511;
+const unsigned NUM_PREDEF_TYPE_IDS = 512;
// Ensure we do not overrun the predefined types we reserved
// in the enum PredefinedTypeIDs above.
diff --git a/clang/lib/CodeGen/CGDebugInfo.cpp b/clang/lib/CodeGen/CGDebugInfo.cpp
index 27bbbfc6f531a1..3f9e14a52fc801 100644
--- a/clang/lib/CodeGen/CGDebugInfo.cpp
+++ b/clang/lib/CodeGen/CGDebugInfo.cpp
@@ -909,6 +909,13 @@ llvm::DIType *CGDebugInfo::CreateType(const BuiltinType *BT) {
TheCU, TheCU->getFile(), 0); \
return SingletonId; \
}
+#define AMDGPU_NAMED_BARRIER_TYPE(Name, Id, SingletonId, Width, Align, Scope) \
+ case BuiltinType::Id: { \
+ if (!SingletonId) \
+ SingletonId = \
+ DBuilder.createBasicType(Name, Width, llvm::dwarf::DW_ATE_unsigned); \
+ return SingletonId;
+ }
#include "clang/Basic/AMDGPUTypes.def"
case BuiltinType::UChar:
case BuiltinType::Char_U:
diff --git a/clang/lib/CodeGen/CodeGenTypes.cpp b/clang/lib/CodeGen/CodeGenTypes.cpp
index f87184fc77832c..09191a4901f493 100644
--- a/clang/lib/CodeGen/CodeGenTypes.cpp
+++ b/clang/lib/CodeGen/CodeGenTypes.cpp
@@ -564,6 +564,10 @@ llvm::Type *CodeGenTypes::ConvertType(QualType T) {
#define AMDGPU_OPAQUE_PTR_TYPE(Name, Id, SingletonId, Width, Align, AS) \
case BuiltinType::Id: \
return llvm::PointerType::get(getLLVMContext(), AS);
+#define AMDGPU_NAMED_BARRIER_TYPE(Name, Id, SingletonId, Width, Align, Scope) \
+ case BuiltinType::Id: \
+ return llvm::TargetExtType::get(getLLVMContext(), "amdgcn.named.barrier", \
+ {}, {Scope});
#include "clang/Basic/AMDGPUTypes.def"
#define HLSL_INTANGIBLE_TYPE(Name, Id, SingletonId) case BuiltinType::Id:
#include "clang/Basic/HLSLIntangibleTypes.def"
diff --git a/clang/test/AST/ast-dump-amdgpu-types.c b/clang/test/AST/ast-dump-amdgpu-types.c
index e032d678f1a09e..f01461cdba2374 100644
--- a/clang/test/AST/ast-dump-amdgpu-types.c
+++ b/clang/test/AST/ast-dump-amdgpu-types.c
@@ -1,10 +1,15 @@
// REQUIRES: amdgpu-registered-target
// Test without serialization:
-// RUN: %clang_cc1 -triple amdgcn -ast-dump -ast-dump-filter __amdgpu_buffer_rsrc_t %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn -ast-dump -ast-dump-filter __amdgpu_buffer_rsrc_t %s | FileCheck %s -check-prefix=BUFFER-RSRC
+// RUN: %clang_cc1 -triple amdgcn -ast-dump -ast-dump-filter __amdgpu_named_workgroup_barrier %s | FileCheck %s -check-prefix=WORKGROUP-BARRIER
//
// Test with serialization:
// RUN: %clang_cc1 -triple amdgcn -emit-pch -o %t %s
-// 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
+// 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
+// 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
-// CHECK: TypedefDecl {{.*}} implicit __amdgpu_buffer_rsrc_t
-// CHECK-NEXT: -BuiltinType {{.*}} '__amdgpu_buffer_rsrc_t'
+// BUFFER-RSRC: TypedefDecl {{.*}} implicit __amdgpu_buffer_rsrc_t
+// BUFFER-RSRC-NEXT: -BuiltinType {{.*}} '__amdgpu_buffer_rsrc_t'
+
+// WORKGROUP-BARRIER: TypedefDecl {{.*}} implicit __amdgpu_named_workgroup_barrier_t
+// WORKGROUP-BARRIER-NEXT: -BuiltinType {{.*}} '__amdgpu_named_workgroup_barrier_t'
diff --git a/clang/test/CodeGen/amdgpu-barrier-type-debug-info.c b/clang/test/CodeGen/amdgpu-barrier-type-debug-info.c
new file mode 100644
index 00000000000000..f595f1b222c4f6
--- /dev/null
+++ b/clang/test/CodeGen/amdgpu-barrier-type-debug-info.c
@@ -0,0 +1,8 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn -emit-llvm -o - %s -debug-info-kind=limited 2>&1 | FileCheck %s
+
+// CHECK: name: "__amdgpu_named_workgroup_barrier_t",{{.*}}baseType: ![[BT:[0-9]+]]
+// CHECK: [[BT]] = !DIBasicType(name: "__amdgpu_named_workgroup_barrier_t", size: 128, encoding: DW_ATE_unsigned)
+void test_locals(void) {
+ __amdgpu_named_workgroup_barrier_t k0;
+}
diff --git a/clang/test/CodeGenCXX/amdgpu-barrier-typeinfo.cpp b/clang/test/CodeGenCXX/amdgpu-barrier-typeinfo.cpp
new file mode 100644
index 00000000000000..a47f217dcd3db6
--- /dev/null
+++ b/clang/test/CodeGenCXX/amdgpu-barrier-typeinfo.cpp
@@ -0,0 +1,10 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn %s -emit-llvm -o - | FileCheck %s
+
+namespace std { class type_info; };
+
+auto &b0 = typeid(__amdgpu_named_workgroup_barrier_t);
+
+// CHECK-DAG: @_ZTSu34__amdgpu_named_workgroup_barrier_t = {{.*}} c"u34__amdgpu_named_workgroup_barrier_t\00"
+// CHECK-DAG: @_ZTIu34__amdgpu_named_workgroup_barrier_t = {{.*}} @_ZTVN10__cxxabiv123__fundamental_type_infoE, {{.*}} @_ZTSu34__amdgpu_named_workgroup_barrier_t
+
diff --git a/clang/test/CodeGenHIP/amdgpu-barrier-type.hip b/clang/test/CodeGenHIP/amdgpu-barrier-type.hip
new file mode 100644
index 00000000000000..229e8b3c737c6a
--- /dev/null
+++ b/clang/test/CodeGenHIP/amdgpu-barrier-type.hip
@@ -0,0 +1,42 @@
+// 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
+
+#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-LABEL: define {{[^@]+}}@_Z7testSemPu34__amdgpu_named_workgroup_barrier_t
+// CHECK-SAME: (ptr noundef [[P:%.*]]) #[[ATTR0:[0-9]+]] {
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[RETVAL:%.*]] = alloca ptr, align 8, addrspace(5)
+// CHECK-NEXT: [[P_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
+// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
+// CHECK-NEXT: [[P_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[P_ADDR]] to ptr
+// 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 ([2 x target("amdgcn.named.barrier", 0)], ptr addrspacecast (ptr addrspace(1) @arr to ptr), i64 0, i64 1)) #[[ATTR2]]
+// 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]]
+// 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) {
+ useBar(p);
+ useBar(&bar);
+ useBar(&arr[1]);
+ useBar(&str.y);
+ useBar(getBar());
+ return getBar();
+}
diff --git a/clang/test/SemaCXX/amdgpu-barrier.cpp b/clang/test/SemaCXX/amdgpu-barrier.cpp
new file mode 100644
index 00000000000000..a171433727dda4
--- /dev/null
+++ b/clang/test/SemaCXX/amdgpu-barrier.cpp
@@ -0,0 +1,17 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -fsyntax-only -verify -std=gnu++11 -triple amdgcn -Wno-unused-value %s
+
+void foo() {
+ int n = 100;
+ __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'}}
+ static_cast<__amdgpu_named_workgroup_barrier_t>(n); // expected-error {{static_cast from 'int' to '__amdgpu_named_workgroup_barrier_t' is not allowed}}
+ 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}}
+ reinterpret_cast<__amdgpu_named_workgroup_barrier_t>(n); // expected-error {{reinterpret_cast from 'int' to '__amdgpu_named_workgroup_barrier_t' is not allowed}}
+ int c(v); // expected-error {{cannot initialize a variable of type 'int' with an lvalue of type '__amdgpu_named_workgroup_barrier_t'}}
+ __amdgpu_named_workgroup_barrier_t k;
+ int *ip = (int *)k; // expected-error {{cannot cast from type '__amdgpu_named_workgroup_barrier_t' to pointer type 'int *'}}
+ void *vp = (void *)k; // expected-error {{cannot cast from type '__amdgpu_named_workgroup_barrier_t' to pointer type 'void *'}}
+}
+
+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
new file mode 100644
index 00000000000000..ccd99b1e2c1f26
--- /dev/null
+++ b/clang/test/SemaHIP/amdgpu-barrier.hip
@@ -0,0 +1,20 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -fsyntax-only -verify -triple amdgcn -Wno-unused-value %s
+// RUN: %clang_cc1 -fsyntax-only -verify -triple x86_64 -aux-triple amdgcn -Wno-unused-value %s
+
+#define __device__ __attribute__((device))
+
+__device__ void foo() {
+ int n = 100;
+ __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'}}
+ static_cast<__amdgpu_named_workgroup_barrier_t>(n); // expected-error {{static_cast from 'int' to '__amdgpu_named_workgroup_barrier_t' is not allowed}}
+ 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}}
+ reinterpret_cast<__amdgpu_named_workgroup_barrier_t>(n); // expected-error {{reinterpret_cast from 'int' to '__amdgpu_named_workgroup_barrier_t' is not allowed}}
+ int c(v); // expected-error {{cannot initialize a variable of type 'int' with an lvalue of type '__amdgpu_named_workgroup_barrier_t'}}
+ __amdgpu_named_workgroup_barrier_t k;
+ int *ip = (int *)k; // expected-error {{cannot cast from type '__amdgpu_named_workgroup_barrier_t' to pointer type 'int *'}}
+ void *vp = (void *)k; // expected-error {{cannot cast from type '__amdgpu_named_workgroup_barrier_t' to pointer type 'void *'}}
+}
+
+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
new file mode 100644
index 00000000000000..150c311c7c5930
--- /dev/null
+++ b/clang/test/SemaOpenCL/amdgpu-barrier.cl
@@ -0,0 +1,12 @@
+// 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
+
+void foo() {
+ 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'}}
+ __amdgpu_named_workgroup_barrier_t k;
+ int *ip = (int *)k; // expected-error {{operand of type '__amdgpu_named_workgroup_barrier_t' where arithmetic or pointer type is required}}
+ void *vp = (void *)k; // expected-error {{operand of type '__amdgpu_named_workgroup_barrier_t' where arithmetic or pointer type is required}}
+ }
diff --git a/clang/test/SemaOpenMP/amdgpu-barrier.cpp b/clang/test/SemaOpenMP/amdgpu-barrier.cpp
new file mode 100644
index 00000000000000..70aaefd080885e
--- /dev/null
+++ b/clang/test/SemaOpenMP/amdgpu-barrier.cpp
@@ -0,0 +1,17 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -Wno-unused-value %s
+
+void foo() {
+#pragma omp target
+ {
+ int n = 100;
+ __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'}}
+ static_cast<__amdgpu_named_workgroup_barrier_t>(n); // expected-error {{static_cast from 'int' to '__amdgpu_named_workgroup_barrier_t' is not allowed}}
+ 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}}
+ reinterpret_cast<__amdgpu_named_workgroup_barrier_t>(n); // expected-error {{reinterpret_cast from 'int' to '__amdgpu_named_workgroup_barrier_t' is not allowed}}
+ int c(v); // expected-error {{cannot initialize a variable of type 'int' with an lvalue of type '__amdgpu_named_workgroup_barrier_t'}}
+ __amdgpu_named_workgroup_barrier_t k;
+ int *ip = (int *)k; // expected-error {{cannot cast from type '__amdgpu_named_workgroup_barrier_t' to pointer type 'int *'}}
+ void *vp = (void *)k; // expected-error {{cannot cast from type '__amdgpu_named_workgroup_barrier_t' to pointer type 'void *'}}
+ }
+ }
diff --git a/llvm/lib/IR/Type.cpp b/llvm/lib/IR/Type.cpp
index f618263f79c313..1c5a97cbc80e19 100644
--- a/llvm/lib/IR/Type.cpp
+++ b/llvm/lib/IR/Type.cpp
@@ -839,6 +839,14 @@ Expected<TargetExtType *> TargetExtType::checkParams(TargetExtType *TTy) {
"target extension type riscv.vector.tuple should have one "
"type parameter and one integer parameter");
+ // Opaque types in the AMDGPU name space.
+ if (TTy->Name == "amdgcn.named.barrier" &&
+ (TTy->getNumTypeParameters() != 0 || TTy->getNumIntParameters() != 1)) {
+ return createStringError("target extension type amdgcn.named.barrier "
+ "should have no type parameters "
+ "and one integer parameter");
+ }
+
return TTy;
}
>From 8212b013170aee371b7554347e9cc49055f918b9 Mon Sep 17 00:00:00 2001
From: gangc <gangc at amd.com>
Date: Thu, 24 Oct 2024 11:18:22 -0700
Subject: [PATCH 2/3] [AMDGPU] Add a type for the named barrier
---
clang/lib/CodeGen/CGDebugInfo.cpp | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/clang/lib/CodeGen/CGDebugInfo.cpp b/clang/lib/CodeGen/CGDebugInfo.cpp
index 3f9e14a52fc801..4d5d2de61c4557 100644
--- a/clang/lib/CodeGen/CGDebugInfo.cpp
+++ b/clang/lib/CodeGen/CGDebugInfo.cpp
@@ -910,11 +910,11 @@ llvm::DIType *CGDebugInfo::CreateType(const BuiltinType *BT) {
return SingletonId; \
}
#define AMDGPU_NAMED_BARRIER_TYPE(Name, Id, SingletonId, Width, Align, Scope) \
- case BuiltinType::Id: { \
+ case BuiltinType::Id: { \
if (!SingletonId) \
SingletonId = \
DBuilder.createBasicType(Name, Width, llvm::dwarf::DW_ATE_unsigned); \
- return SingletonId;
+ return SingletonId; \
}
#include "clang/Basic/AMDGPUTypes.def"
case BuiltinType::UChar:
>From fccd6f6f44ecc0c0165b441e3edf47834bf07d12 Mon Sep 17 00:00:00 2001
From: gangc <gangc at amd.com>
Date: Thu, 24 Oct 2024 16:10:48 -0700
Subject: [PATCH 3/3] [AMDGPU] add the type for named barrier
opaque type in getTargetTypeInfo
---
llvm/lib/IR/Type.cpp | 6 ++++++
1 file changed, 6 insertions(+)
diff --git a/llvm/lib/IR/Type.cpp b/llvm/lib/IR/Type.cpp
index 1c5a97cbc80e19..9365c4fa276d9e 100644
--- a/llvm/lib/IR/Type.cpp
+++ b/llvm/lib/IR/Type.cpp
@@ -892,6 +892,12 @@ static TargetTypeInfo getTargetTypeInfo(const TargetExtType *Ty) {
if (Name.starts_with("dx."))
return TargetTypeInfo(PointerType::get(C, 0));
+ // Opaque types in the AMDGPU name space.
+ if (Name == "amdgcn.named.barrier") {
+ return TargetTypeInfo(FixedVectorType::get(Type::getInt32Ty(C), 4),
+ TargetExtType::CanBeGlobal);
+ }
+
return TargetTypeInfo(Type::getVoidTy(C));
}
More information about the cfe-commits
mailing list