https://github.com/ayokunle321 updated https://github.com/llvm/llvm-project/pull/223137
>From 945988b5bdf85770d9a9e4251b4d5c364474cd83 Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Sat, 12 Sep 2026 13:02:02 +0200 Subject: [PATCH 1/2] [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 8d23951dd64ba9..7a52f41d074084 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -1044,34 +1044,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 00000000000000..ce8bef760069e9 --- /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 00000000000000..a0c2918fcdcc48 --- /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 00000000000000..8db9ac72bf963e --- /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); +} + >From ee3a8636c0af3c912ad60dc315ed9809d0dbf885 Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Sat, 10 Oct 2026 01:03:58 +0200 Subject: [PATCH 2/2] add arg types to checks --- .../builtins-amdgcn-raw-buffer-atomic-add.hip | 16 ++++++++-------- .../builtins-amdgcn-raw-buffer-atomic-max.hip | 8 ++++---- 2 files changed, 12 insertions(+), 12 deletions(-) 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 index a0c2918fcdcc48..a401393b257e1c 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-buffer-atomic-add.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-buffer-atomic-add.hip @@ -57,36 +57,36 @@ 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: 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 {{.*}}) +// 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: 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 {{.*}}) +// 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: 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 {{.*}}) +// 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: 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> {{.*}}) +// 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 index 8db9ac72bf963e..2b248731d3f5b1 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-buffer-atomic-max.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-raw-buffer-atomic-max.hip @@ -57,18 +57,18 @@ 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 +// 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 {{.*}}) +// 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: 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 {{.*}}) +// 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
