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 <ga...@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 <ga...@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 <ga...@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)); } _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits