[AMDGPU] Add a type for the named barrier #113614
Merged
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
llvmbot
added
clang
|
@llvm/pr-subscribers-clang-codegen @llvm/pr-subscribers-debuginfo Author: Gang Chen (cmc-rep) ChangesFull diff: https://github.com/llvm/llvm-project/pull/113614.diff 13 Files Affected:
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;
}
|
|
@llvm/pr-subscribers-backend-amdgpu Author: Gang Chen (cmc-rep) ChangesFull diff: https://github.com/llvm/llvm-project/pull/113614.diff 13 Files Affected:
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;
}
|
|
@llvm/pr-subscribers-clang-modules Author: Gang Chen (cmc-rep) ChangesFull diff: https://github.com/llvm/llvm-project/pull/113614.diff 13 Files Affected:
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;
}
|
|
@llvm/pr-subscribers-clang Author: Gang Chen (cmc-rep) ChangesFull diff: https://github.com/llvm/llvm-project/pull/113614.diff 13 Files Affected:
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;
}
|
|
✅ With the latest revision this PR passed the C/C++ code formatter. |
opaque type in getTargetTypeInfo
jayfoad
approved these changes
Oct 25, 2024
fix newline fix no-external-type-id.cppm
accidentally delete an empty space line in test
yxsamliu
approved these changes
Oct 25, 2024
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/144/builds/10134 Here is the relevant piece of the build log for the reference |
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/162/builds/9088 Here is the relevant piece of the build log for the reference |
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/89/builds/9137 Here is the relevant piece of the build log for the reference |
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/2/builds/9625 Here is the relevant piece of the build log for the reference |
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/78/builds/8414 Here is the relevant piece of the build log for the reference |
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/30/builds/8924 Here is the relevant piece of the build log for the reference |
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/157/builds/11082 Here is the relevant piece of the build log for the reference |
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/153/builds/12808 Here is the relevant piece of the build log for the reference |
|
Hi, |
|
I just pushed a fix in 75252e2. |
|
How did this got through pre-merge testing, I wonder. |
|
This PR did bump |
|
Ah, so it's the issue of PRs not being tested against |
I see another error in one of our bots for this patch: https://lab.llvm.org/staging/#/builders/130/builds/7112 |
|
Looks like I also need to fix the number in that test |
|
fix to the test no-external-type-id.cppm: #113738 |
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/60/builds/11107 Here is the relevant piece of the build log for the reference |
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/56/builds/10701 Here is the relevant piece of the build log for the reference |
jayfoad
reviewed
Oct 29, 2024
| @@ -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)) { | |||
There was a problem hiding this comment.
This could also be a good place to check that the "scope" argument is 0, if that is the only supported value.
NoumanAmir657
pushed a commit
to NoumanAmir657/llvm-project
that referenced
this pull request
Nov 4, 2024
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
No description provided.