https://github.com/ranapratap55 updated
https://github.com/llvm/llvm-project/pull/160258
>From 2da220874eec49becbbdbfc0a346228d8e71e1a1 Mon Sep 17 00:00:00 2001
From: ranapratap55 <ranapratapreddy.nimmakay...@amd.com>
Date: Tue, 23 Sep 2025 14:46:16 +0530
Subject: [PATCH 1/3] [AMDGPU] Add a new builtin type for image descriptor rsrc
---
clang/include/clang/Basic/AMDGPUTypes.def | 7 +++++++
clang/include/clang/Basic/Builtins.def | 1 +
clang/lib/AST/ASTContext.cpp | 4 ++++
clang/lib/CodeGen/CGDebugInfo.cpp | 8 +++++++
clang/lib/CodeGen/CodeGenTypes.cpp | 4 ++++
.../amdgpu-image-rsrc-type-debug-info.c | 18 ++++++++++++++++
.../CodeGenCXX/amdgpu-image-rsrc-typeinfo.cpp | 7 +++++++
clang/test/SemaCXX/amdgpu-image-rsrc.cpp | 21 +++++++++++++++++++
clang/test/SemaOpenCL/amdgpu-image-rsrc.cl | 14 +++++++++++++
clang/test/SemaOpenMP/amdgpu-image-rsrc.cpp | 12 +++++++++++
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 1 +
11 files changed, 97 insertions(+)
create mode 100644 clang/test/CodeGen/amdgpu-image-rsrc-type-debug-info.c
create mode 100644 clang/test/CodeGenCXX/amdgpu-image-rsrc-typeinfo.cpp
create mode 100644 clang/test/SemaCXX/amdgpu-image-rsrc.cpp
create mode 100644 clang/test/SemaOpenCL/amdgpu-image-rsrc.cl
create mode 100644 clang/test/SemaOpenMP/amdgpu-image-rsrc.cpp
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 12c7d48e20d67..c2f0534f5ffe6 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"], []>;
>From 2ceded10f785c66fb3758621e40dd88b31c3214a Mon Sep 17 00:00:00 2001
From: ranapratap55 <ranapratapreddy.nimmakay...@amd.com>
Date: Wed, 24 Sep 2025 10:37:12 +0530
Subject: [PATCH 2/3] [AMDGPU] Change image desc rsrc character from 'c' to 't'
---
clang/include/clang/Basic/Builtins.def | 2 +-
clang/lib/AST/ASTContext.cpp | 2 +-
2 files changed, 2 insertions(+), 2 deletions(-)
diff --git a/clang/include/clang/Basic/Builtins.def
b/clang/include/clang/Basic/Builtins.def
index a91315680f93f..f621a72dfa92b 100644
--- a/clang/include/clang/Basic/Builtins.def
+++ b/clang/include/clang/Basic/Builtins.def
@@ -34,7 +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.
+// Qt -> 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 7ba1dfed1c0db..c60c53720b908 100644
--- a/clang/lib/AST/ASTContext.cpp
+++ b/clang/lib/AST/ASTContext.cpp
@@ -12580,7 +12580,7 @@ static QualType DecodeTypeFromStr(const char *&Str,
const ASTContext &Context,
Type = Context.AMDGPUBufferRsrcTy;
break;
}
- case 'c': {
+ case 't': {
Type = Context.AMDGPUImageDescRsrcTy;
break;
}
>From ef4abb91e1dca51e047f6d15a7784e3c1fe57de2 Mon Sep 17 00:00:00 2001
From: ranapratap55 <ranapratapreddy.nimmakay...@amd.com>
Date: Thu, 25 Sep 2025 11:26:18 +0530
Subject: [PATCH 3/3] [AMDGPU] Removed def in intrinsic(not related to this
clang builtin patch)
---
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 1 -
1 file changed, 1 deletion(-)
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index d41ce5b64b7cd..afce1fe6af854 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -967,7 +967,6 @@ 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"], []>;
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
