[clang] [CIR][AMDGPU] Add support for AMDGCN raw buffer store format builtins (PR #223112)
Ayokunle Amodu via cfe-commits
cfe-commits at lists.llvm.org
Fri Oct 9 15:54:01 PDT 2026
https://github.com/ayokunle321 updated https://github.com/llvm/llvm-project/pull/223112
>From 90b83f36abde758bfe5346a253fbcd0cb2d3673f Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Sat, 12 Sep 2026 03:11:07 +0200
Subject: [PATCH 1/3] 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 8d23951dd64ba9..92991ac1c79b80 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -1033,6 +1033,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 00000000000000..52c658d7275e2d
--- /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);
+}
+
>From 592c7ebfb474f1ae6fa39a72d4d7803721f70cbf Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Thu, 17 Sep 2026 15:09:56 +0200
Subject: [PATCH 2/3] add raw buffer store
---
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 6 ++++
...uiltins-amdgcn-raw-buffer-store-format.hip | 35 +++++++++++++++++++
2 files changed, 41 insertions(+)
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index 92991ac1c79b80..717ef4c1d83510 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -1039,6 +1039,12 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
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
index 52c658d7275e2d..774541863f5b30 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-buffer-store-format.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-buffer-store-format.hip
@@ -52,3 +52,38 @@ __device__ void test_raw_buffer_store_format_v4f32_non_const_soffset(v4f32 vdata
__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({{.*}})
+__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({{.*}})
+__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({{.*}})
+__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({{.*}})
+__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);
+}
>From 01b6013d83bdf667fb748ee8df1f5d0b998a7b85 Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Sat, 10 Oct 2026 00:53:34 +0200
Subject: [PATCH 3/3] add args to llvm checks
---
.../builtins-amdgcn-raw-buffer-store-format.hip | 16 ++++++++--------
1 file changed, 8 insertions(+), 8 deletions(-)
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
index 774541863f5b30..e41e04b85a2336 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-buffer-store-format.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-buffer-store-format.hip
@@ -19,7 +19,7 @@ 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> {{.*}})
+// 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);
@@ -28,7 +28,7 @@ __device__ void test_raw_buffer_store_format_v4f32(v4f32 vdata, __amdgpu_buffer_
// 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> {{.*}})
+// 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);
@@ -37,7 +37,7 @@ __device__ void test_raw_buffer_store_format_v4f16(v4f16 vdata, __amdgpu_buffer_
// 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> {{.*}})
+// 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);
@@ -46,7 +46,7 @@ __device__ void test_raw_buffer_store_format_v4f32_non_const_offset(v4f32 vdata,
// 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> {{.*}})
+// 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);
@@ -55,7 +55,7 @@ __device__ void test_raw_buffer_store_format_v4f32_non_const_soffset(v4f32 vdata
// 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({{.*}})
+// 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);
@@ -64,7 +64,7 @@ __device__ v4f32 test_raw_buffer_load_format_v4f32(__amdgpu_buffer_rsrc_t rsrc,
// 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({{.*}})
+// 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);
@@ -73,7 +73,7 @@ __device__ v4f16 test_raw_buffer_load_format_v4f16(__amdgpu_buffer_rsrc_t rsrc,
// 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({{.*}})
+// 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);
@@ -82,7 +82,7 @@ __device__ v4f32 test_raw_buffer_load_format_v4f32_non_const_offset(__amdgpu_buf
// 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({{.*}})
+// 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);
More information about the cfe-commits
mailing list