[clang] [CIR][AMDGPU] Adds `__amdgpu_buffer_rsrc_t` in the buffer-resource address space (PR #204782)
Rana Pratap Reddy via cfe-commits
cfe-commits at lists.llvm.org
Mon Jul 6 01:51:50 PDT 2026
https://github.com/ranapratap55 updated https://github.com/llvm/llvm-project/pull/204782
>From bdb31e942a1b563f49a708ce9b5633b65e4515ee Mon Sep 17 00:00:00 2001
From: ranapratap55 <RanaPratapReddy.Nimmakayala at amd.com>
Date: Fri, 19 Jun 2026 14:02:32 +0530
Subject: [PATCH 1/2] [CIR][AMDGPU] Adds __amdgpu_buffer_rsrc_t in the
buffer-resource address space
---
clang/lib/CIR/CodeGen/CIRGenTypes.cpp | 4 +-
.../builtins-amdgcn-buffer-rsrc-type.hip | 80 +++++++++++++++++++
2 files changed, 83 insertions(+), 1 deletion(-)
create mode 100644 clang/test/CIR/CodeGenHIP/builtins-amdgcn-buffer-rsrc-type.hip
diff --git a/clang/lib/CIR/CodeGen/CIRGenTypes.cpp b/clang/lib/CIR/CodeGen/CIRGenTypes.cpp
index 3170666304a06..6c12728fecea1 100644
--- a/clang/lib/CIR/CodeGen/CIRGenTypes.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenTypes.cpp
@@ -500,7 +500,9 @@ mlir::Type CIRGenTypes::convertType(QualType type) {
if (BuiltinType::Id == BuiltinType::AMDGPUTexture) { \
resultType = cir::VectorType::get(builder.getSInt32Ty(), 8); \
} else { \
- resultType = builder.getPointerTo(cgm.voidTy); \
+ resultType = builder.getPointerTo( \
+ cgm.voidTy, \
+ cir::TargetAddressSpaceAttr::get(&getMLIRContext(), AS)); \
} \
break; \
}
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-buffer-rsrc-type.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-buffer-rsrc-type.hip
new file mode 100644
index 0000000000000..84fa7d9f74c3b
--- /dev/null
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-buffer-rsrc-type.hip
@@ -0,0 +1,80 @@
+#include "../CodeGenCUDA/Inputs/cuda.h"
+
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
+// RUN: -target-cpu gfx1100 -fcuda-is-device -emit-cir %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
+// RUN: -target-cpu gfx1100 -fcuda-is-device -emit-llvm %s -o %t-cir.ll
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t-cir.ll %s
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \
+// RUN: -target-cpu gfx1100 -fcuda-is-device -emit-llvm %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
+
+struct BufferResourceHolder {
+ int x;
+ __amdgpu_buffer_rsrc_t r;
+};
+
+__device__ void consume_buffer(__amdgpu_buffer_rsrc_t);
+__device__ __amdgpu_buffer_rsrc_t make_resource();
+
+// CIR-LABEL: cir.func {{.*}}test_buffer_rsrc_passthrough
+// CIR-SAME: !cir.ptr<!void, target_address_space(8)>
+// CIR-SAME: -> !cir.ptr<!void, target_address_space(8)>
+// LLVM-LABEL: define{{.*}} ptr addrspace(8) @{{.*}}test_buffer_rsrc_passthrough
+__device__ __amdgpu_buffer_rsrc_t
+test_buffer_rsrc_passthrough(__amdgpu_buffer_rsrc_t rsrc) {
+ return rsrc;
+}
+
+// CIR-LABEL: cir.func {{.*}}test_buffer_rsrc_load
+// CIR: cir.load {{.*}} : !cir.ptr<!cir.ptr<!void, target_address_space(8)>>, !cir.ptr<!void, target_address_space(8)>
+// LLVM-LABEL: define{{.*}} ptr addrspace(8) @{{.*}}test_buffer_rsrc_load
+__device__ __amdgpu_buffer_rsrc_t
+test_buffer_rsrc_load(__amdgpu_buffer_rsrc_t *p) {
+ return *p;
+}
+
+// CIR-LABEL: cir.func {{.*}}test_buffer_rsrc_store
+// CIR: cir.store{{.*}} : !cir.ptr<!void, target_address_space(8)>,
+// LLVM-LABEL: define{{.*}}@{{.*}}test_buffer_rsrc_store
+__device__ void
+test_buffer_rsrc_store(__amdgpu_buffer_rsrc_t *p, __amdgpu_buffer_rsrc_t rsrc) {
+ *p = rsrc;
+}
+
+// CIR-LABEL: cir.func {{.*}}test_struct_member
+// CIR: cir.get_member {{.*}} {name = "r"} {{.*}} -> !cir.ptr<!cir.ptr<!void, target_address_space(8)>>
+// CIR: cir.load {{.*}} : !cir.ptr<!cir.ptr<!void, target_address_space(8)>>, !cir.ptr<!void, target_address_space(8)>
+// LLVM-LABEL: define{{.*}} ptr addrspace(8) @{{.*}}test_struct_member
+__device__ __amdgpu_buffer_rsrc_t test_struct_member(BufferResourceHolder *a) {
+ return a->r;
+}
+
+// CIR-LABEL: cir.func {{.*}}test_pass_by_value
+// CIR: cir.call {{.*}}consume_buffer{{.*}}!cir.ptr<!void, target_address_space(8)>
+// LLVM-LABEL: define{{.*}}@{{.*}}test_pass_by_value
+// LLVM: call void @{{.*}}consume_buffer{{.*}}(ptr addrspace(8)
+__device__ void test_pass_by_value(__amdgpu_buffer_rsrc_t rsrc) {
+ consume_buffer(rsrc);
+}
+
+// CIR-LABEL: cir.func {{.*}}test_call_returns_resource
+// CIR: cir.call {{.*}}make_resource{{.*}} -> !cir.ptr<!void, target_address_space(8)>
+// LLVM-LABEL: define{{.*}} ptr addrspace(8) @{{.*}}test_call_returns_resource
+__device__ __amdgpu_buffer_rsrc_t test_call_returns_resource() {
+ return make_resource();
+}
+
+// CIR-LABEL: cir.func {{.*}}test_return_struct
+// CIR: cir.get_member {{.*}} {name = "r"} {{.*}} -> !cir.ptr<!cir.ptr<!void, target_address_space(8)>>
+// LLVM-LABEL: define{{.*}}@{{.*}}test_return_struct
+__device__ BufferResourceHolder test_return_struct(__amdgpu_buffer_rsrc_t rsrc) {
+ BufferResourceHolder a;
+ a.x = 0;
+ a.r = rsrc;
+ return a;
+}
>From 11d3fba4587e29b7933c1473f1f4adee5822b662 Mon Sep 17 00:00:00 2001
From: ranapratap55 <RanaPratapReddy.Nimmakayala at amd.com>
Date: Mon, 6 Jul 2026 14:21:33 +0530
Subject: [PATCH 2/2] [fix] rename test file name
---
...ns-amdgcn-buffer-rsrc-type.hip => amdgcn-buffer-rsrc-type.hip} | 0
1 file changed, 0 insertions(+), 0 deletions(-)
rename clang/test/CIR/CodeGenHIP/{builtins-amdgcn-buffer-rsrc-type.hip => amdgcn-buffer-rsrc-type.hip} (100%)
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-buffer-rsrc-type.hip b/clang/test/CIR/CodeGenHIP/amdgcn-buffer-rsrc-type.hip
similarity index 100%
rename from clang/test/CIR/CodeGenHIP/builtins-amdgcn-buffer-rsrc-type.hip
rename to clang/test/CIR/CodeGenHIP/amdgcn-buffer-rsrc-type.hip
More information about the cfe-commits
mailing list