llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-debuginfo Author: Rana Pratap Reddy (ranapratap55) <details> <summary>Changes</summary> Adding a new builtin type for AMDGPU's image descriptor rsrc data type This requires for https://github.com/llvm/llvm-project/pull/140210 --- Full diff: https://github.com/llvm/llvm-project/pull/160258.diff 11 Files Affected: - (modified) clang/include/clang/Basic/AMDGPUTypes.def (+7) - (modified) clang/include/clang/Basic/Builtins.def (+1) - (modified) clang/lib/AST/ASTContext.cpp (+4) - (modified) clang/lib/CodeGen/CGDebugInfo.cpp (+8) - (modified) clang/lib/CodeGen/CodeGenTypes.cpp (+4) - (added) clang/test/CodeGen/amdgpu-image-rsrc-type-debug-info.c (+18) - (added) clang/test/CodeGenCXX/amdgpu-image-rsrc-typeinfo.cpp (+7) - (added) clang/test/SemaCXX/amdgpu-image-rsrc.cpp (+21) - (added) clang/test/SemaOpenCL/amdgpu-image-rsrc.cl (+14) - (added) clang/test/SemaOpenMP/amdgpu-image-rsrc.cpp (+12) - (modified) llvm/include/llvm/IR/IntrinsicsAMDGPU.td (+1) ``````````diff diff --git a/clang/include/clang/Basic/AMDGPUTypes.def b/clang/include/clang/Basic/AMDGPUTypes.def index d3dff446f9edf..8c0bd73252c50 100644 --- a/clang/include/clang/Basic/AMDGPUTypes.def +++ b/clang/include/clang/Basic/AMDGPUTypes.def @@ -20,10 +20,17 @@ AMDGPU_TYPE(Name, Id, SingletonId, Width, Align) #endif +#ifndef AMDGPU_IMAGE_RSRC_TYPE +#define AMDGPU_IMAGE_RSRC_TYPE(Name, Id, SingletonId) \ + AMDGPU_TYPE(Name, Id, SingletonId, 256, 256) +#endif + AMDGPU_OPAQUE_PTR_TYPE("__amdgpu_buffer_rsrc_t", AMDGPUBufferRsrc, AMDGPUBufferRsrcTy, 128, 128, 8) +AMDGPU_IMAGE_RSRC_TYPE("__amdgpu_image_rsrc_t", AMDGPUImageDescRsrc, AMDGPUImageDescRsrcTy) 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 +#undef AMDGPU_IMAGE_RSRC_TYPE diff --git a/clang/include/clang/Basic/Builtins.def b/clang/include/clang/Basic/Builtins.def index 48437c9397570..a91315680f93f 100644 --- a/clang/include/clang/Basic/Builtins.def +++ b/clang/include/clang/Basic/Builtins.def @@ -34,6 +34,7 @@ // Q -> target builtin type, followed by a character to distinguish the builtin type // Qa -> AArch64 svcount_t builtin type. // Qb -> AMDGPU __amdgpu_buffer_rsrc_t builtin type. +// Qc -> AMDGPU __amdgpu_image_desc_t builtin type. // E -> ext_vector, followed by the number of elements and the base type. // X -> _Complex, followed by the base type. // Y -> ptrdiff_t diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp index 97c59b2ceec2f..7ba1dfed1c0db 100644 --- a/clang/lib/AST/ASTContext.cpp +++ b/clang/lib/AST/ASTContext.cpp @@ -12580,6 +12580,10 @@ static QualType DecodeTypeFromStr(const char *&Str, const ASTContext &Context, Type = Context.AMDGPUBufferRsrcTy; break; } + case 'c': { + Type = Context.AMDGPUImageDescRsrcTy; + break; + } default: llvm_unreachable("Unexpected target builtin type"); } diff --git a/clang/lib/CodeGen/CGDebugInfo.cpp b/clang/lib/CodeGen/CGDebugInfo.cpp index 578d09f7971d6..3f5277ab66f55 100644 --- a/clang/lib/CodeGen/CGDebugInfo.cpp +++ b/clang/lib/CodeGen/CGDebugInfo.cpp @@ -1020,6 +1020,14 @@ llvm::DIType *CGDebugInfo::CreateType(const BuiltinType *BT) { DBuilder.createBasicType(Name, Width, llvm::dwarf::DW_ATE_unsigned); \ return SingletonId; \ } +#define AMDGPU_IMAGE_RSRC_TYPE(Name, Id, SingletonId) \ + case BuiltinType::Id: { \ + if (!SingletonId) \ + SingletonId = \ + DBuilder.createForwardDecl(llvm::dwarf::DW_TAG_structure_type, Name, \ + TheCU, TheCU->getFile(), 0); \ + 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 3ffe999d01178..e3e44556ce514 100644 --- a/clang/lib/CodeGen/CodeGenTypes.cpp +++ b/clang/lib/CodeGen/CodeGenTypes.cpp @@ -581,6 +581,10 @@ llvm::Type *CodeGenTypes::ConvertType(QualType T) { case BuiltinType::Id: \ return llvm::TargetExtType::get(getLLVMContext(), "amdgcn.named.barrier", \ {}, {Scope}); +#define AMDGPU_IMAGE_RSRC_TYPE(Name, Id, SingletonId) \ + case BuiltinType::Id: \ + return llvm::VectorType::get(llvm::Type::getInt32Ty(getLLVMContext()), 8, \ + false); #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/CodeGen/amdgpu-image-rsrc-type-debug-info.c b/clang/test/CodeGen/amdgpu-image-rsrc-type-debug-info.c new file mode 100644 index 0000000000000..0e42420e26322 --- /dev/null +++ b/clang/test/CodeGen/amdgpu-image-rsrc-type-debug-info.c @@ -0,0 +1,18 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5 +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn -emit-llvm -o - %s -debug-info-kind=limited | FileCheck %s + +// CHECK-LABEL: define dso_local void @test_locals( +// CHECK-SAME: ) #[[ATTR0:[0-9]+]] !dbg [[DBG6:![0-9]+]] { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: [[IMG:%.*]] = alloca <8 x i32>, align 32, addrspace(5) +// CHECK-NEXT: [[IMG_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[IMG]] to ptr +// CHECK-NEXT: #dbg_declare(ptr addrspace(5) [[IMG]], [[META11:![0-9]+]], !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef), [[META14:![0-9]+]]) +// CHECK-NEXT: [[TMP0:%.*]] = load <8 x i32>, ptr [[IMG_ASCAST]], align 32, !dbg [[DBG15:![0-9]+]] +// CHECK-NEXT: ret void, !dbg [[DBG16:![0-9]+]] +// +void test_locals(void) { + __amdgpu_image_rsrc_t img; + (void)img; +} + diff --git a/clang/test/CodeGenCXX/amdgpu-image-rsrc-typeinfo.cpp b/clang/test/CodeGenCXX/amdgpu-image-rsrc-typeinfo.cpp new file mode 100644 index 0000000000000..d96cf5f35c5b2 --- /dev/null +++ b/clang/test/CodeGenCXX/amdgpu-image-rsrc-typeinfo.cpp @@ -0,0 +1,7 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5 +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn %s -emit-llvm -o - | FileCheck %s +namespace std { class type_info; } +auto &a = typeid(__amdgpu_image_rsrc_t); +//// NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line: +// CHECK: {{.*}} diff --git a/clang/test/SemaCXX/amdgpu-image-rsrc.cpp b/clang/test/SemaCXX/amdgpu-image-rsrc.cpp new file mode 100644 index 0000000000000..1a19a94039b5e --- /dev/null +++ b/clang/test/SemaCXX/amdgpu-image-rsrc.cpp @@ -0,0 +1,21 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5 +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -fsyntax-only -verify -std=gnu++11 -triple amdgcn -Wno-unused-value %s + +void foo() { + int n = 1; + __amdgpu_image_rsrc_t v = 0; // expected-error {{cannot initialize a variable of type '__amdgpu_image_rsrc_t' with an rvalue of type 'int'}} + static_cast<__amdgpu_image_rsrc_t>(n); // expected-error {{static_cast from 'int' to '__amdgpu_image_rsrc_t' is not allowed}} + reinterpret_cast<__amdgpu_image_rsrc_t>(n); // expected-error {{reinterpret_cast from 'int' to '__amdgpu_image_rsrc_t' is not allowed}} + (void)(v + v); // expected-error {{invalid operands}} + int x(v); // expected-error {{cannot initialize a variable of type 'int' with an lvalue of type '__amdgpu_image_rsrc_t'}} + __amdgpu_image_rsrc_t k; +} + +static_assert(sizeof(__amdgpu_image_rsrc_t) == 32, "size"); +static_assert(alignof(__amdgpu_image_rsrc_t) == 32, "align"); + +template<class T> void bar(T); +void use(__amdgpu_image_rsrc_t r) { bar(r); } +struct S { __amdgpu_image_rsrc_t r; int a; }; +static_assert(sizeof(S) == 64, "struct layout"); diff --git a/clang/test/SemaOpenCL/amdgpu-image-rsrc.cl b/clang/test/SemaOpenCL/amdgpu-image-rsrc.cl new file mode 100644 index 0000000000000..341ab667ebd06 --- /dev/null +++ b/clang/test/SemaOpenCL/amdgpu-image-rsrc.cl @@ -0,0 +1,14 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5 +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -verify -cl-std=CL1.2 -triple amdgcn-amd-amdhsa %s +// RUN: %clang_cc1 -verify -cl-std=CL2.0 -triple amdgcn-amd-amdhsa %s + +void f() { + int n = 3; + __amdgpu_image_rsrc_t v = 0; // expected-error {{initializing '__private __amdgpu_image_rsrc_t' with an expression of incompatible type 'int'}} + int k = v; // expected-error {{initializing '__private int' with an expression of incompatible type '__private __amdgpu_image_rsrc_t'}} + (void)(v + v); // expected-error {{invalid operands}} + __amdgpu_image_rsrc_t r; + int *p = (int*)r; // expected-error {{operand of type '__amdgpu_image_rsrc_t' where arithmetic or pointer type is required}} + (void)p; +} diff --git a/clang/test/SemaOpenMP/amdgpu-image-rsrc.cpp b/clang/test/SemaOpenMP/amdgpu-image-rsrc.cpp new file mode 100644 index 0000000000000..91d566be9b8a3 --- /dev/null +++ b/clang/test/SemaOpenMP/amdgpu-image-rsrc.cpp @@ -0,0 +1,12 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5 +// 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 = 5; + __amdgpu_image_rsrc_t v = 0; // expected-error {{cannot initialize a variable of type '__amdgpu_image_rsrc_t' with an rvalue of type 'int'}} + (void)(v + v); // expected-error {{invalid operands to binary expression ('__amdgpu_image_rsrc_t' and '__amdgpu_image_rsrc_t'}} + } +} diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td index afce1fe6af854..d41ce5b64b7cd 100644 --- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td +++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td @@ -967,6 +967,7 @@ class AMDGPUDimProps<bits<3> enc, string name, string asmsuffix, bits<8> NumGradients = !size(GradientArgs); } +def AMDGPUImageDescRsrcTy : LLVMType<v8i32>; def AMDGPUDim1D : AMDGPUDimProps<0x0, "1d", "1D", ["s"], []>; def AMDGPUDim2D : AMDGPUDimProps<0x1, "2d", "2D", ["s", "t"], []>; def AMDGPUDim3D : AMDGPUDimProps<0x2, "3d", "3D", ["s", "t", "r"], []>; `````````` </details> https://github.com/llvm/llvm-project/pull/160258 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
