https://github.com/ayokunle321 created https://github.com/llvm/llvm-project/pull/223112
Adds codegen for the following AMDGCN raw buffer store format builtins: - __builtin_amdgcn_raw_buffer_store_format_v4f32 (v4f32) - __builtin_amdgcn_raw_buffer_store_format_v4f16 (v4f16) These are lowered to the corresponding `llvm.amdgcn.raw.ptr.buffer.store.format` intrinsic. >From 857288c6d11b4ced55282539a6d2e745619f0650 Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Sat, 12 Sep 2026 03:11:07 +0200 Subject: [PATCH] add raw buffer store codegen --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 6 +++ ...uiltins-amdgcn-raw-buffer-store-format.hip | 54 +++++++++++++++++++ 2 files changed, 60 insertions(+) create mode 100644 clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-buffer-store-format.hip diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 1c23ea142bf8d..4f7911dd24359 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -1011,6 +1011,12 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, getContext().BuiltinInfo.getName(builtinId)); return mlir::Value{}; } + case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_format_v4f32: + case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_format_v4f16: + return emitBuiltinWithOneOverloadedType<5>( + expr, "amdgcn.raw.ptr.buffer.store.format", + cir::VoidType::get(builder.getContext())) + .getValue(); case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b8: case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b16: case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b32: diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-buffer-store-format.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-buffer-store-format.hip new file mode 100644 index 0000000000000..52c658d7275e2 --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-buffer-store-format.hip @@ -0,0 +1,54 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgpu8.03-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -fcuda-is-device -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s + +// RUN: %clang_cc1 -triple amdgpu8.03-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgpu8.03-amd-amdhsa -x hip -std=c++11 \ +// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +#define __device__ __attribute__((device)) + +typedef float v4f32 __attribute__((ext_vector_type(4))); +typedef _Float16 v4f16 __attribute__((ext_vector_type(4))); + +// CIR-LABEL: @_Z34test_raw_buffer_store_format_v4f32Dv4_fu22__amdgpu_buffer_rsrc_tii +// CIR: cir.call_llvm_intrinsic "amdgcn.raw.ptr.buffer.store.format" {{.*}} : (!cir.vector<4 x !cir.float>, !cir.ptr<!void, target_address_space(8)>, !s32i, !s32i, !s32i) -> !void +// LLVM: define{{.*}} void @_Z34test_raw_buffer_store_format_v4f32Dv4_fu22__amdgpu_buffer_rsrc_tii +// LLVM: call{{.*}} void @llvm.amdgcn.raw.ptr.buffer.store.format.v4f32(<4 x float> {{.*}}) +__device__ void test_raw_buffer_store_format_v4f32(v4f32 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, + int soffset) { + __builtin_amdgcn_raw_buffer_store_format_v4f32(vdata, rsrc, 0, 0, 0); +} + +// CIR-LABEL: @_Z34test_raw_buffer_store_format_v4f16Dv4_DF16_u22__amdgpu_buffer_rsrc_tii +// CIR: cir.call_llvm_intrinsic "amdgcn.raw.ptr.buffer.store.format" {{.*}} : (!cir.vector<4 x !cir.f16>, !cir.ptr<!void, target_address_space(8)>, !s32i, !s32i, !s32i) -> !void +// LLVM: define{{.*}} void @_Z34test_raw_buffer_store_format_v4f16Dv4_DF16_u22__amdgpu_buffer_rsrc_tii +// LLVM: call{{.*}} void @llvm.amdgcn.raw.ptr.buffer.store.format.v4f16(<4 x half> {{.*}}) +__device__ void test_raw_buffer_store_format_v4f16(v4f16 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, + int soffset) { + __builtin_amdgcn_raw_buffer_store_format_v4f16(vdata, rsrc, 0, 0, 0); +} + +// CIR-LABEL: @_Z51test_raw_buffer_store_format_v4f32_non_const_offsetDv4_fu22__amdgpu_buffer_rsrc_tii +// CIR: cir.call_llvm_intrinsic "amdgcn.raw.ptr.buffer.store.format" {{.*}} : (!cir.vector<4 x !cir.float>, !cir.ptr<!void, target_address_space(8)>, !s32i, !s32i, !s32i) -> !void +// LLVM: define{{.*}} void @_Z51test_raw_buffer_store_format_v4f32_non_const_offsetDv4_fu22__amdgpu_buffer_rsrc_tii +// LLVM: call{{.*}} void @llvm.amdgcn.raw.ptr.buffer.store.format.v4f32(<4 x float> {{.*}}) +__device__ void test_raw_buffer_store_format_v4f32_non_const_offset(v4f32 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, + int soffset) { + __builtin_amdgcn_raw_buffer_store_format_v4f32(vdata, rsrc, offset, 0, 0); +} + +// CIR-LABEL: @_Z52test_raw_buffer_store_format_v4f32_non_const_soffsetDv4_fu22__amdgpu_buffer_rsrc_tii +// CIR: cir.call_llvm_intrinsic "amdgcn.raw.ptr.buffer.store.format" {{.*}} : (!cir.vector<4 x !cir.float>, !cir.ptr<!void, target_address_space(8)>, !s32i, !s32i, !s32i) -> !void +// LLVM: define{{.*}} void @_Z52test_raw_buffer_store_format_v4f32_non_const_soffsetDv4_fu22__amdgpu_buffer_rsrc_tii +// LLVM: call{{.*}} void @llvm.amdgcn.raw.ptr.buffer.store.format.v4f32(<4 x float> {{.*}}) +__device__ void test_raw_buffer_store_format_v4f32_non_const_soffset(v4f32 vdata, __amdgpu_buffer_rsrc_t rsrc, int offset, + int soffset) { + __builtin_amdgcn_raw_buffer_store_format_v4f32(vdata, rsrc, 0, soffset, 0); +} + _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
