https://github.com/ranapratap55 updated https://github.com/llvm/llvm-project/pull/204782
>From bdb31e942a1b563f49a708ce9b5633b65e4515ee Mon Sep 17 00:00:00 2001 From: ranapratap55 <[email protected]> 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 <[email protected]> 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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
