Author: Ayokunle Amodu Date: 2026-10-12T06:01:55-04:00 New Revision: 71e68dd7b851b1f5e05f9614d7d26002fb7d1dc7
URL: https://github.com/llvm/llvm-project/commit/71e68dd7b851b1f5e05f9614d7d26002fb7d1dc7 DIFF: https://github.com/llvm/llvm-project/commit/71e68dd7b851b1f5e05f9614d7d26002fb7d1dc7.diff LOG: [CIR][AMDGPU] Add support for raw ptr buffer atomic builtins (#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 --------- Co-authored-by: Claude Opus 5 <[email protected]> Added: clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx13.hip clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-buffer-atomic-add.hip clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-buffer-atomic-max.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 a9db1082c097f..7e5c743e3d699 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -1116,34 +1116,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: return emitBuiltinWithOneOverloadedType<2>( expr, "amdgcn.s.prefetch.data", 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..2da82e521271e --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx13.hip @@ -0,0 +1,40 @@ +// 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 diff erence +// between CIR (no flags) and classic codegen ('contract'). + +#define __device__ __attribute__((device)) +#define global __attribute__((address_space(1))) + +// 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.ptr<!void, target_address_space(8)>, !s32i, !s32i, !s32i) -> !cir.float +// CIR: cir.call_llvm_intrinsic "amdgcn.raw.ptr.buffer.atomic.fmax" {{.*}} : (!cir.float, !cir.ptr<!void, target_address_space(8)>, !s32i, !s32i, !s32i) -> !cir.float +// CIR: cir.call_llvm_intrinsic "amdgcn.raw.ptr.buffer.atomic.fmin" {{.*}} : (!cir.double, !cir.ptr<!void, target_address_space(8)>, !s32i, !s32i, !s32i) -> !cir.double +// CIR: cir.call_llvm_intrinsic "amdgcn.raw.ptr.buffer.atomic.fmax" {{.*}} : (!cir.double, !cir.ptr<!void, target_address_space(8)>, !s32i, !s32i, !s32i) -> !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 %{{.*}}, ptr addrspace(8) %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 0) +// LLVM: call{{.*}} float @llvm.amdgcn.raw.ptr.buffer.atomic.fmax.f32(float %{{.*}}, ptr addrspace(8) %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 0) +// LLVM: call{{.*}} double @llvm.amdgcn.raw.ptr.buffer.atomic.fmin.f64(double %{{.*}}, ptr addrspace(8) %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 0) +// LLVM: call{{.*}} double @llvm.amdgcn.raw.ptr.buffer.atomic.fmax.f64(double %{{.*}}, ptr addrspace(8) %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 0) +__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..a8db02d4322fb --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-buffer-atomic-add.hip @@ -0,0 +1,93 @@ +// 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 diff erence +// between CIR (no flags) and classic codegen ('contract'). + +#define __device__ __attribute__((device)) + +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, !cir.ptr<!void, target_address_space(8)>, !s32i, !s32i, !s32i) -> !s32i +// LLVM: define{{.*}} i32 @_Z19test_atomic_add_i32u22__amdgpu_buffer_rsrc_tiii +// LLVM: call{{.*}} i32 @llvm.amdgcn.raw.ptr.buffer.atomic.add.i32(i32 %{{.*}}, ptr addrspace(8) %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 0) +__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, !cir.ptr<!void, target_address_space(8)>, !s32i, !s32i, !s32i) -> !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 %{{.*}}, ptr addrspace(8) %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 0) +__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, !cir.ptr<!void, target_address_space(8)>, !s32i, !s32i, !s32i) -> !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 %{{.*}}, ptr addrspace(8) %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 0) +__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>, !cir.ptr<!void, target_address_space(8)>, !s32i, !s32i, !s32i) -> !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> %{{.*}}, ptr addrspace(8) %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 0) +__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..e7e32c04df58a --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-buffer-atomic-max.hip @@ -0,0 +1,73 @@ +// 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 diff erence +// between CIR (no flags) and classic codegen ('contract'). + +#define __device__ __attribute__((device)) + +// CIR-LABEL: @_Z20test_atomic_fmax_f32u22__amdgpu_buffer_rsrc_tfii +// CIR: cir.call_llvm_intrinsic "amdgcn.raw.ptr.buffer.atomic.fmax" {{.*}} : (!cir.float, !cir.ptr<!void, target_address_space(8)>, !s32i, !s32i, !s32i) -> !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 %{{.*}}, ptr addrspace(8) %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 0) +__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, !cir.ptr<!void, target_address_space(8)>, !s32i, !s32i, !s32i) -> !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 %{{.*}}, ptr addrspace(8) %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 0) +__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
