[clang] [CIR][AMDGPU] Add support for raw ptr buffer atomic builtins (PR #223137)
via cfe-commits
cfe-commits at lists.llvm.org
Sat Sep 12 04:06:41 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang
@llvm/pr-subscribers-backend-amdgpu
Author: Ayokunle Amodu (ayokunle321)
<details>
<summary>Changes</summary>
Adds codegen for the following AMDGCN raw pointer buffer atomic builtins:
- __builtin_amdgcn_raw_ptr_buffer_atomic_add_i32
- __builtin_amdgcn_raw_ptr_buffer_atomic_fadd_{f32,f64,v2f16}
- __builtin_amdgcn_raw_ptr_buffer_atomic_fmin_{f32,f64}
- __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_{f32,f64}
These are lowered to the corresponding `llvm.amdgcn.raw.ptr.buffer.atomic.*` intrinsics.
Co-Authored-By: Claude Opus 5
---
Full diff: https://github.com/llvm/llvm-project/pull/223137.diff
4 Files Affected:
- (modified) clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp (+16-24)
- (added) clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx13.hip (+42)
- (added) clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-buffer-atomic-add.hip (+94)
- (added) clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-buffer-atomic-max.hip (+76)
``````````diff
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index 1c23ea142bf8d..f4fa815f93bb0 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -1022,34 +1022,26 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
getContext().BuiltinInfo.getName(builtinId));
return mlir::Value{};
}
- case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_add_i32: {
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented AMDGPU builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
- }
+ case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_add_i32:
+ return emitBuiltinWithOneOverloadedType<5>(
+ expr, "amdgcn.raw.ptr.buffer.atomic.add")
+ .getValue();
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f32:
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f64:
- case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16: {
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented AMDGPU builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
- }
+ case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16:
+ return emitBuiltinWithOneOverloadedType<5>(
+ expr, "amdgcn.raw.ptr.buffer.atomic.fadd")
+ .getValue();
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f32:
- case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f64: {
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented AMDGPU builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
- }
+ case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f64:
+ return emitBuiltinWithOneOverloadedType<5>(
+ expr, "amdgcn.raw.ptr.buffer.atomic.fmin")
+ .getValue();
case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32:
- case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64: {
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented AMDGPU builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
- }
+ case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64:
+ return emitBuiltinWithOneOverloadedType<5>(
+ expr, "amdgcn.raw.ptr.buffer.atomic.fmax")
+ .getValue();
case AMDGPU::BI__builtin_amdgcn_s_prefetch_data: {
cgm.errorNYI(expr->getSourceRange(),
std::string("unimplemented AMDGPU builtin call: ") +
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx13.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx13.hip
new file mode 100644
index 0000000000000..ce8bef760069e
--- /dev/null
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx13.hip
@@ -0,0 +1,42 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgpu13.10-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 amdgpu13.10-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 amdgpu13.10-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
+
+// FIXME: CIR doesn't propagate the 'contract' fast-math flag to LLVM IR
+// calls yet, so LLVM check lines use {{.*}} to tolerate the difference
+// between CIR (no flags) and classic codegen ('contract').
+
+#define __device__ __attribute__((device))
+#define global __attribute__((address_space(1)))
+
+typedef _Float16 v2f16 __attribute__((ext_vector_type(2)));
+
+// CIR-LABEL: @_Z36test_raw_ptr_buffer_atomic_fmin_fmaxPU3AS1fPU3AS1du22__amdgpu_buffer_rsrc_tfdii
+// CIR: cir.call_llvm_intrinsic "amdgcn.raw.ptr.buffer.atomic.fmin" {{.*}} -> !cir.float
+// CIR: cir.call_llvm_intrinsic "amdgcn.raw.ptr.buffer.atomic.fmax" {{.*}} -> !cir.float
+// CIR: cir.call_llvm_intrinsic "amdgcn.raw.ptr.buffer.atomic.fmin" {{.*}} -> !cir.double
+// CIR: cir.call_llvm_intrinsic "amdgcn.raw.ptr.buffer.atomic.fmax" {{.*}} -> !cir.double
+// LLVM: define{{.*}} void @_Z36test_raw_ptr_buffer_atomic_fmin_fmaxPU3AS1fPU3AS1du22__amdgpu_buffer_rsrc_tfdii
+// LLVM: call{{.*}} float @llvm.amdgcn.raw.ptr.buffer.atomic.fmin.f32(float {{.*}})
+// LLVM: call{{.*}} float @llvm.amdgcn.raw.ptr.buffer.atomic.fmax.f32(float {{.*}})
+// LLVM: call{{.*}} double @llvm.amdgcn.raw.ptr.buffer.atomic.fmin.f64(double {{.*}})
+// LLVM: call{{.*}} double @llvm.amdgcn.raw.ptr.buffer.atomic.fmax.f64(double {{.*}})
+__device__ void test_raw_ptr_buffer_atomic_fmin_fmax(global float *fout,
+ global double *dout,
+ __amdgpu_buffer_rsrc_t rsrc,
+ float f, double d, int offset,
+ int soffset) {
+ *fout = __builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f32(f, rsrc, offset, soffset, 0);
+ *fout = __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32(f, rsrc, offset, soffset, 0);
+ *dout = __builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f64(d, rsrc, offset, soffset, 0);
+ *dout = __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64(d, rsrc, offset, soffset, 0);
+}
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-buffer-atomic-add.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-buffer-atomic-add.hip
new file mode 100644
index 0000000000000..a0c2918fcdcc4
--- /dev/null
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-buffer-atomic-add.hip
@@ -0,0 +1,94 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgpu9.0a-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 amdgpu9.42-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 amdgpu9.50-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 amdgpu12.50-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 amdgpu9.0a-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 amdgpu9.42-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 amdgpu9.50-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 amdgpu12.50-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 amdgpu9.0a-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
+
+// RUN: %clang_cc1 -triple amdgpu9.42-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
+
+// RUN: %clang_cc1 -triple amdgpu9.50-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
+
+// RUN: %clang_cc1 -triple amdgpu12.50-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
+
+// FIXME: CIR doesn't propagate the 'contract' fast-math flag to LLVM IR
+// calls yet, so LLVM check lines use {{.*}} to tolerate the difference
+// between CIR (no flags) and classic codegen ('contract').
+
+#define __device__ __attribute__((device))
+#define global __attribute__((address_space(1)))
+
+typedef _Float16 v2f16 __attribute__((ext_vector_type(2)));
+
+// CIR-LABEL: @_Z19test_atomic_add_i32u22__amdgpu_buffer_rsrc_tiii
+// CIR: cir.call_llvm_intrinsic "amdgcn.raw.ptr.buffer.atomic.add" {{.*}} -> !s32i
+// LLVM: define{{.*}} i32 @_Z19test_atomic_add_i32u22__amdgpu_buffer_rsrc_tiii
+// LLVM: call{{.*}} i32 @llvm.amdgcn.raw.ptr.buffer.atomic.add.i32(i32 {{.*}})
+__device__ int test_atomic_add_i32(__amdgpu_buffer_rsrc_t rsrc, int x, int offset,
+ int soffset) {
+ return __builtin_amdgcn_raw_ptr_buffer_atomic_add_i32(x, rsrc, offset, soffset, 0);
+}
+
+// CIR-LABEL: @_Z20test_atomic_fadd_f32u22__amdgpu_buffer_rsrc_tfii
+// CIR: cir.call_llvm_intrinsic "amdgcn.raw.ptr.buffer.atomic.fadd" {{.*}} -> !cir.float
+// LLVM: define{{.*}} float @_Z20test_atomic_fadd_f32u22__amdgpu_buffer_rsrc_tfii
+// LLVM: call{{.*}} float @llvm.amdgcn.raw.ptr.buffer.atomic.fadd.f32(float {{.*}})
+__device__ float test_atomic_fadd_f32(__amdgpu_buffer_rsrc_t rsrc, float x, int offset,
+ int soffset) {
+ return __builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f32(x, rsrc, offset, soffset, 0);
+}
+
+// CIR-LABEL: @_Z20test_atomic_fadd_f64u22__amdgpu_buffer_rsrc_tdii
+// CIR: cir.call_llvm_intrinsic "amdgcn.raw.ptr.buffer.atomic.fadd" {{.*}} -> !cir.double
+// LLVM: define{{.*}} double @_Z20test_atomic_fadd_f64u22__amdgpu_buffer_rsrc_tdii
+// LLVM: call{{.*}} double @llvm.amdgcn.raw.ptr.buffer.atomic.fadd.f64(double {{.*}})
+__device__ double test_atomic_fadd_f64(__amdgpu_buffer_rsrc_t rsrc, double x, int offset,
+ int soffset) {
+ return __builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f64(x, rsrc, offset, soffset, 0);
+}
+
+// CIR-LABEL: @_Z22test_atomic_fadd_v2f16u22__amdgpu_buffer_rsrc_tDv2_DF16_ii
+// CIR: cir.call_llvm_intrinsic "amdgcn.raw.ptr.buffer.atomic.fadd" {{.*}} -> !cir.vector<2 x !cir.f16>
+// LLVM: define{{.*}} <2 x half> @_Z22test_atomic_fadd_v2f16u22__amdgpu_buffer_rsrc_tDv2_DF16_ii
+// LLVM: call{{.*}} <2 x half> @llvm.amdgcn.raw.ptr.buffer.atomic.fadd.v2f16(<2 x half> {{.*}})
+__device__ v2f16 test_atomic_fadd_v2f16(__amdgpu_buffer_rsrc_t rsrc, v2f16 x, int offset,
+ int soffset) {
+ return __builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16(x, rsrc, offset, soffset, 0);
+}
+
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-buffer-atomic-max.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-buffer-atomic-max.hip
new file mode 100644
index 0000000000000..8db9ac72bf963
--- /dev/null
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-buffer-atomic-max.hip
@@ -0,0 +1,76 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgpu10.1-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 amdgpu10.3-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 amdgpu12.50-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 amdgpu13.10-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 amdgpu10.1-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 amdgpu10.3-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 amdgpu12.50-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 amdgpu13.10-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 amdgpu10.1-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
+
+// RUN: %clang_cc1 -triple amdgpu10.3-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
+
+// RUN: %clang_cc1 -triple amdgpu12.50-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
+
+// RUN: %clang_cc1 -triple amdgpu13.10-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
+
+// FIXME: CIR doesn't propagate the 'contract' fast-math flag to LLVM IR
+// calls yet, so LLVM check lines use {{.*}} to tolerate the difference
+// between CIR (no flags) and classic codegen ('contract').
+
+#define __device__ __attribute__((device))
+#define global __attribute__((address_space(1)))
+
+typedef _Float16 v2f16 __attribute__((ext_vector_type(2)));
+
+// CIR-LABEL: @_Z20test_atomic_fmax_f32u22__amdgpu_buffer_rsrc_tfii
+// CIR: cir.call_llvm_intrinsic "amdgcn.raw.ptr.buffer.atomic.fmax" {{.*}} -> !cir.float
+// LLVM: define{{.*}} float @_Z20test_atomic_fmax_f32u22__amdgpu_buffer_rsrc_tfii
+// LLVM: call{{.*}} float @llvm.amdgcn.raw.ptr.buffer.atomic.fmax.f32(float {{.*}})
+__device__ float test_atomic_fmax_f32(__amdgpu_buffer_rsrc_t rsrc, float x, int offset,
+ int soffset) {
+ return __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32(x, rsrc, offset, soffset, 0);
+}
+
+// CIR-LABEL: @_Z20test_atomic_fmax_f64u22__amdgpu_buffer_rsrc_tdii
+// CIR: cir.call_llvm_intrinsic "amdgcn.raw.ptr.buffer.atomic.fmax" {{.*}} -> !cir.double
+// LLVM: define{{.*}} double @_Z20test_atomic_fmax_f64u22__amdgpu_buffer_rsrc_tdii
+// LLVM: call{{.*}} double @llvm.amdgcn.raw.ptr.buffer.atomic.fmax.f64(double {{.*}})
+__device__ double test_atomic_fmax_f64(__amdgpu_buffer_rsrc_t rsrc, double x, int offset,
+ int soffset) {
+ return __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64(x, rsrc, offset, soffset, 0);
+}
+
``````````
</details>
https://github.com/llvm/llvm-project/pull/223137
More information about the cfe-commits
mailing list