[clang] 533bab2 - [CIR][AMDGPU] Add support for AMDGCN raw buffer store format builtins (#223112)

[via cfe-commits]

Author: Ayokunle Amodu
Date: 2026-10-12T06:02:29-04:00
New Revision: 533bab2e8a43c69c39b911e0fa3f37addaaf45f5

URL: https://github.com/llvm/llvm-project/commit/533bab2e8a43c69c39b911e0fa3f37addaaf45f5
DIFF: https://github.com/llvm/llvm-project/commit/533bab2e8a43c69c39b911e0fa3f37addaaf45f5.diff

LOG: [CIR][AMDGPU] Add support for AMDGCN raw buffer store format builtins (#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.

Added: 
    clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-buffer-store-format.hip

Modified: 
    clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp

Removed: 
    


################################################################################
diff  --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index 7e5c743e3d699..275974b924760 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -1105,6 +1105,18 @@ 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_format_v4f32:
+  case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_format_v4f16:
+    return emitBuiltinWithOneOverloadedType<4>(
+               expr, "amdgcn.raw.ptr.buffer.load.format",
+               convertType(expr->getType()))
+        .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..e41e04b85a233
--- /dev/null
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-buffer-store-format.hip
@@ -0,0 +1,89 @@
+// 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> %{{.*}}, ptr addrspace(8) %{{.*}}, i32 0, i32 0, i32 0)
+__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> %{{.*}}, ptr addrspace(8) %{{.*}}, i32 0, i32 0, i32 0)
+__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> %{{.*}}, ptr addrspace(8) %{{.*}}, i32 %{{.*}}, i32 0, i32 0)
+__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> %{{.*}}, ptr addrspace(8) %{{.*}}, i32 0, i32 %{{.*}}, i32 0)
+__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);
+}
+
+// CIR-LABEL: @_Z33test_raw_buffer_load_format_v4f32u22__amdgpu_buffer_rsrc_tii
+// CIR: cir.call_llvm_intrinsic "amdgcn.raw.ptr.buffer.load.format" {{.*}} : (!cir.ptr<!void, target_address_space(8)>, !s32i, !s32i, !s32i) -> !cir.vector<4 x !cir.float>
+// LLVM: define{{.*}} <4 x float> @_Z33test_raw_buffer_load_format_v4f32u22__amdgpu_buffer_rsrc_tii
+// LLVM: call{{.*}} <4 x float> @llvm.amdgcn.raw.ptr.buffer.load.format.v4f32(ptr addrspace(8) %{{.*}}, i32 0, i32 0, i32 0)
+__device__ v4f32 test_raw_buffer_load_format_v4f32(__amdgpu_buffer_rsrc_t rsrc, int offset,
+                                                     int soffset) {
+  return __builtin_amdgcn_raw_buffer_load_format_v4f32(rsrc, 0, 0, 0);
+}
+
+// CIR-LABEL: @_Z33test_raw_buffer_load_format_v4f16u22__amdgpu_buffer_rsrc_tii
+// CIR: cir.call_llvm_intrinsic "amdgcn.raw.ptr.buffer.load.format" {{.*}} : (!cir.ptr<!void, target_address_space(8)>, !s32i, !s32i, !s32i) -> !cir.vector<4 x !cir.f16>
+// LLVM: define{{.*}} <4 x half> @_Z33test_raw_buffer_load_format_v4f16u22__amdgpu_buffer_rsrc_tii
+// LLVM: call{{.*}} <4 x half> @llvm.amdgcn.raw.ptr.buffer.load.format.v4f16(ptr addrspace(8) %{{.*}}, i32 0, i32 0, i32 0)
+__device__ v4f16 test_raw_buffer_load_format_v4f16(__amdgpu_buffer_rsrc_t rsrc, int offset,
+                                                     int soffset) {
+  return __builtin_amdgcn_raw_buffer_load_format_v4f16(rsrc, 0, 0, 0);
+}
+
+// CIR-LABEL: @_Z50test_raw_buffer_load_format_v4f32_non_const_offsetu22__amdgpu_buffer_rsrc_tii
+// CIR: cir.call_llvm_intrinsic "amdgcn.raw.ptr.buffer.load.format" {{.*}} : (!cir.ptr<!void, target_address_space(8)>, !s32i, !s32i, !s32i) -> !cir.vector<4 x !cir.float>
+// LLVM: define{{.*}} <4 x float> @_Z50test_raw_buffer_load_format_v4f32_non_const_offsetu22__amdgpu_buffer_rsrc_tii
+// LLVM: call{{.*}} <4 x float> @llvm.amdgcn.raw.ptr.buffer.load.format.v4f32(ptr addrspace(8) %{{.*}}, i32 %{{.*}}, i32 0, i32 0)
+__device__ v4f32 test_raw_buffer_load_format_v4f32_non_const_offset(__amdgpu_buffer_rsrc_t rsrc, int offset,
+                                                                      int soffset) {
+  return __builtin_amdgcn_raw_buffer_load_format_v4f32(rsrc, offset, 0, 0);
+}
+
+// CIR-LABEL: @_Z51test_raw_buffer_load_format_v4f32_non_const_soffsetu22__amdgpu_buffer_rsrc_tii
+// CIR: cir.call_llvm_intrinsic "amdgcn.raw.ptr.buffer.load.format" {{.*}} : (!cir.ptr<!void, target_address_space(8)>, !s32i, !s32i, !s32i) -> !cir.vector<4 x !cir.float>
+// LLVM: define{{.*}} <4 x float> @_Z51test_raw_buffer_load_format_v4f32_non_const_soffsetu22__amdgpu_buffer_rsrc_tii
+// LLVM: call{{.*}} <4 x float> @llvm.amdgcn.raw.ptr.buffer.load.format.v4f32(ptr addrspace(8) %{{.*}}, i32 0, i32 %{{.*}}, i32 0)
+__device__ v4f32 test_raw_buffer_load_format_v4f32_non_const_soffset(__amdgpu_buffer_rsrc_t rsrc, int offset,
+                                                                       int soffset) {
+  return __builtin_amdgcn_raw_buffer_load_format_v4f32(rsrc, 0, soffset, 0);
+}