https://github.com/ayokunle321 created
https://github.com/llvm/llvm-project/pull/223137
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
>From be85498589fd3f2275f3b08b52c9e0e22a6b6e11 Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <[email protected]>
Date: Sat, 12 Sep 2026 13:02:02 +0200
Subject: [PATCH] [CIR][AMDGPU] Add support for raw ptr buffer atomic builtins
Adds codegen for the following AMDGCN raw pointer buffer atomic builtins.
These are lowered to the corresponding llvm.amdgcn.raw.ptr.buffer.atomic.*
intrinsics.
- `__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}`
Each intrinsic is overloaded on its first operand and returns that same type,
so the result type does not need to be passed to
emitBuiltinWithOneOverloadedType.
The tests mirror the OG test files and their target sets: fmin is only covered
by builtins-amdgcn-gfx13.cl, so it is tested there rather than alongside fmax.
Co-Authored-By: Claude Opus 5 <[email protected]>
---
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 40 ++++----
.../CIR/CodeGenHIP/builtins-amdgcn-gfx13.hip | 42 +++++++++
.../builtins-amdgcn-raw-buffer-atomic-add.hip | 94 +++++++++++++++++++
.../builtins-amdgcn-raw-buffer-atomic-max.hip | 76 +++++++++++++++
4 files changed, 228 insertions(+), 24 deletions(-)
create mode 100644 clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx13.hip
create mode 100644
clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-buffer-atomic-add.hip
create mode 100644
clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-buffer-atomic-max.hip
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);
+}
+
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
