https://github.com/arsenm updated https://github.com/llvm/llvm-project/pull/213119
>From 13a1fcb3a06c2120432d200e7026f01df606fe39 Mon Sep 17 00:00:00 2001 From: Matt Arsenault <[email protected]> Date: Thu, 30 Jul 2026 21:44:57 +0200 Subject: [PATCH] AMDGPU: Add missing atomic-fmin-fmax-global to gfx13 feature map fillAMDGCNFeatureMap omitted atomic-fmin-fmax-global-f32 and atomic-fmin-fmax-global-f64 for gfx1310/gfx13-generic, so clang wrongly rejected the raw_ptr_buffer_atomic_f{min,max}_f{32,64} builtins on those targets even though the backend enables the features. Add them to the gfx13 case. Co-authored-by: Claude (Claude-Opus-4.8) --- .../CodeGenOpenCL/builtins-amdgcn-gfx13.cl | 54 ++++++++ .../builtins-amdgcn-raw-buffer-atomic-max.cl | 1 + llvm/lib/TargetParser/AMDGPUTargetParser.cpp | 2 + .../AMDGPU/fp-min-max-buffer-ptr-atomics.ll | 122 ++++++++++++++++++ .../AMDGPU/fp64-min-max-buffer-ptr-atomics.ll | 116 +++++++++++++++++ 5 files changed, 295 insertions(+) diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl index a6e0b15921860..af40e044e2099 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl @@ -46,3 +46,57 @@ void test_cvt_scalef32_pk32_fp6_f32(global uint6 *out, float32 srcf32, float src *out = __builtin_amdgcn_cvt_scalef32_pk32_fp6_f32(srcf32, src); } +// CHECK-LABEL: @test_raw_ptr_buffer_atomic_fmin_fmax( +// CHECK-NEXT: entry: +// CHECK-NEXT: [[FOUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5) +// CHECK-NEXT: [[DOUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5) +// CHECK-NEXT: [[RSRC_ADDR:%.*]] = alloca ptr addrspace(8), align 16, addrspace(5) +// CHECK-NEXT: [[F_ADDR:%.*]] = alloca float, align 4, addrspace(5) +// CHECK-NEXT: [[D_ADDR:%.*]] = alloca double, align 8, addrspace(5) +// CHECK-NEXT: [[OFFSET_ADDR:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-NEXT: [[SOFFSET_ADDR:%.*]] = alloca i32, align 4, addrspace(5) +// CHECK-NEXT: store ptr addrspace(1) [[FOUT:%.*]], ptr addrspace(5) [[FOUT_ADDR]], align 8 +// CHECK-NEXT: store ptr addrspace(1) [[DOUT:%.*]], ptr addrspace(5) [[DOUT_ADDR]], align 8 +// CHECK-NEXT: store ptr addrspace(8) [[RSRC:%.*]], ptr addrspace(5) [[RSRC_ADDR]], align 16 +// CHECK-NEXT: store float [[F:%.*]], ptr addrspace(5) [[F_ADDR]], align 4 +// CHECK-NEXT: store double [[D:%.*]], ptr addrspace(5) [[D_ADDR]], align 8 +// CHECK-NEXT: store i32 [[OFFSET:%.*]], ptr addrspace(5) [[OFFSET_ADDR]], align 4 +// CHECK-NEXT: store i32 [[SOFFSET:%.*]], ptr addrspace(5) [[SOFFSET_ADDR]], align 4 +// CHECK-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(5) [[F_ADDR]], align 4 +// CHECK-NEXT: [[TMP1:%.*]] = load ptr addrspace(8), ptr addrspace(5) [[RSRC_ADDR]], align 16 +// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(5) [[OFFSET_ADDR]], align 4 +// CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr addrspace(5) [[SOFFSET_ADDR]], align 4 +// CHECK-NEXT: [[TMP4:%.*]] = call float @llvm.amdgcn.raw.ptr.buffer.atomic.fmin.f32(float [[TMP0]], ptr addrspace(8) [[TMP1]], i32 [[TMP2]], i32 [[TMP3]], i32 0) +// CHECK-NEXT: [[TMP5:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[FOUT_ADDR]], align 8 +// CHECK-NEXT: store float [[TMP4]], ptr addrspace(1) [[TMP5]], align 4 +// CHECK-NEXT: [[TMP6:%.*]] = load float, ptr addrspace(5) [[F_ADDR]], align 4 +// CHECK-NEXT: [[TMP7:%.*]] = load ptr addrspace(8), ptr addrspace(5) [[RSRC_ADDR]], align 16 +// CHECK-NEXT: [[TMP8:%.*]] = load i32, ptr addrspace(5) [[OFFSET_ADDR]], align 4 +// CHECK-NEXT: [[TMP9:%.*]] = load i32, ptr addrspace(5) [[SOFFSET_ADDR]], align 4 +// CHECK-NEXT: [[TMP10:%.*]] = call float @llvm.amdgcn.raw.ptr.buffer.atomic.fmax.f32(float [[TMP6]], ptr addrspace(8) [[TMP7]], i32 [[TMP8]], i32 [[TMP9]], i32 0) +// CHECK-NEXT: [[TMP11:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[FOUT_ADDR]], align 8 +// CHECK-NEXT: store float [[TMP10]], ptr addrspace(1) [[TMP11]], align 4 +// CHECK-NEXT: [[TMP12:%.*]] = load double, ptr addrspace(5) [[D_ADDR]], align 8 +// CHECK-NEXT: [[TMP13:%.*]] = load ptr addrspace(8), ptr addrspace(5) [[RSRC_ADDR]], align 16 +// CHECK-NEXT: [[TMP14:%.*]] = load i32, ptr addrspace(5) [[OFFSET_ADDR]], align 4 +// CHECK-NEXT: [[TMP15:%.*]] = load i32, ptr addrspace(5) [[SOFFSET_ADDR]], align 4 +// CHECK-NEXT: [[TMP16:%.*]] = call double @llvm.amdgcn.raw.ptr.buffer.atomic.fmin.f64(double [[TMP12]], ptr addrspace(8) [[TMP13]], i32 [[TMP14]], i32 [[TMP15]], i32 0) +// CHECK-NEXT: [[TMP17:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[DOUT_ADDR]], align 8 +// CHECK-NEXT: store double [[TMP16]], ptr addrspace(1) [[TMP17]], align 8 +// CHECK-NEXT: [[TMP18:%.*]] = load double, ptr addrspace(5) [[D_ADDR]], align 8 +// CHECK-NEXT: [[TMP19:%.*]] = load ptr addrspace(8), ptr addrspace(5) [[RSRC_ADDR]], align 16 +// CHECK-NEXT: [[TMP20:%.*]] = load i32, ptr addrspace(5) [[OFFSET_ADDR]], align 4 +// CHECK-NEXT: [[TMP21:%.*]] = load i32, ptr addrspace(5) [[SOFFSET_ADDR]], align 4 +// CHECK-NEXT: [[TMP22:%.*]] = call double @llvm.amdgcn.raw.ptr.buffer.atomic.fmax.f64(double [[TMP18]], ptr addrspace(8) [[TMP19]], i32 [[TMP20]], i32 [[TMP21]], i32 0) +// CHECK-NEXT: [[TMP23:%.*]] = load ptr addrspace(1), ptr addrspace(5) [[DOUT_ADDR]], align 8 +// CHECK-NEXT: store double [[TMP22]], ptr addrspace(1) [[TMP23]], align 8 +// CHECK-NEXT: ret void +// +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/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-max.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-max.cl index 6928032d8c45a..feb8682d522b1 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-max.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-atomic-max.cl @@ -2,6 +2,7 @@ // RUN: %clang_cc1 -triple amdgpu10.1-unknown-unknown -emit-llvm -o - %s | FileCheck %s // RUN: %clang_cc1 -triple amdgpu10.3-unknown-unknown -emit-llvm -o - %s | FileCheck %s // RUN: %clang_cc1 -triple amdgpu12.50-unknown-unknown -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgpu13.10-unknown-unknown -emit-llvm -o - %s | FileCheck %s // REQUIRES: amdgpu-registered-target // CHECK-LABEL: define dso_local float @test_atomic_fmax_f32( diff --git a/llvm/lib/TargetParser/AMDGPUTargetParser.cpp b/llvm/lib/TargetParser/AMDGPUTargetParser.cpp index c55e81bbdf6ff..a244b1acecc53 100644 --- a/llvm/lib/TargetParser/AMDGPUTargetParser.cpp +++ b/llvm/lib/TargetParser/AMDGPUTargetParser.cpp @@ -517,6 +517,8 @@ static void fillAMDGCNFeatureMap(StringRef GPU, const Triple &T, Features["atomic-flat-pk-add-16-insts"] = true; Features["atomic-global-pk-add-bf16-inst"] = true; Features["atomic-ds-pk-add-16-insts"] = true; + Features["atomic-fmin-fmax-global-f32"] = true; + Features["atomic-fmin-fmax-global-f64"] = true; Features["s-wakeup-barrier-inst"] = true; Features["f16bf16-to-fp6bf6-cvt-scale-insts"] = true; Features["f32-to-fp6bf6-cvt-scale-insts"] = true; diff --git a/llvm/test/CodeGen/AMDGPU/fp-min-max-buffer-ptr-atomics.ll b/llvm/test/CodeGen/AMDGPU/fp-min-max-buffer-ptr-atomics.ll index 6e5ec14404a28..768b1d5cfdbda 100644 --- a/llvm/test/CodeGen/AMDGPU/fp-min-max-buffer-ptr-atomics.ll +++ b/llvm/test/CodeGen/AMDGPU/fp-min-max-buffer-ptr-atomics.ll @@ -4,12 +4,14 @@ ; RUN: llc < %s -mtriple=amdgpu10.10 | FileCheck %s -check-prefix=GFX10 ; RUN: llc < %s -mtriple=amdgpu10.30 | FileCheck %s -check-prefix=GFX1030 ; RUN: llc < %s -mtriple=amdgpu11.00 | FileCheck %s -check-prefix=GFX1100 +; RUN: llc < %s -mtriple=amdgpu13.10 | FileCheck %s -check-prefix=GFX1310 ; RUN: llc < %s -global-isel -mtriple=amdgpu6.01 | FileCheck %s -check-prefix=G_SI ; RUN: llc < %s -global-isel -mtriple=amdgpu7.01 | FileCheck %s -check-prefix=G_GFX7 ; RUN: llc < %s -global-isel -mtriple=amdgpu10.10 | FileCheck %s -check-prefix=G_GFX10 ; RUN: llc < %s -global-isel -mtriple=amdgpu10.30 | FileCheck %s -check-prefix=G_GFX1030 ; RUN: llc < %s -global-isel -mtriple=amdgpu11.00 | FileCheck %s -check-prefix=G_GFX1100 +; RUN: llc < %s -global-isel -mtriple=amdgpu13.10 | FileCheck %s -check-prefix=G_GFX1310 declare float @llvm.amdgcn.raw.ptr.buffer.atomic.fmin.f32(float, ptr addrspace(8), i32, i32, i32 immarg) declare float @llvm.amdgcn.raw.ptr.buffer.atomic.fmax.f32(float, ptr addrspace(8), i32, i32, i32 immarg) @@ -68,6 +70,16 @@ define amdgpu_kernel void @raw_ptr_buffer_atomic_min_noret_f32(ptr addrspace(8) ; GFX1100-NEXT: buffer_atomic_min_f32 v0, v1, s[0:3], 0 offen ; GFX1100-NEXT: s_endpgm ; +; GFX1310-LABEL: raw_ptr_buffer_atomic_min_noret_f32: +; GFX1310: ; %bb.0: ; %main_body +; GFX1310-NEXT: s_clause 0x1 +; GFX1310-NEXT: s_load_b64 s[6:7], s[4:5], 0x34 nv +; GFX1310-NEXT: s_load_b128 s[0:3], s[4:5], 0x24 nv +; GFX1310-NEXT: s_wait_kmcnt 0x0 +; GFX1310-NEXT: v_dual_mov_b32 v0, s6 :: v_dual_mov_b32 v1, s7 +; GFX1310-NEXT: buffer_atomic_min_num_f32 v0, v1, s[0:3], null offen +; GFX1310-NEXT: s_endpgm +; ; G_SI-LABEL: raw_ptr_buffer_atomic_min_noret_f32: ; G_SI: ; %bb.0: ; %main_body ; G_SI-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0xd @@ -119,6 +131,16 @@ define amdgpu_kernel void @raw_ptr_buffer_atomic_min_noret_f32(ptr addrspace(8) ; G_GFX1100-NEXT: v_dual_mov_b32 v0, s6 :: v_dual_mov_b32 v1, s7 ; G_GFX1100-NEXT: buffer_atomic_min_f32 v0, v1, s[0:3], 0 offen ; G_GFX1100-NEXT: s_endpgm +; +; G_GFX1310-LABEL: raw_ptr_buffer_atomic_min_noret_f32: +; G_GFX1310: ; %bb.0: ; %main_body +; G_GFX1310-NEXT: s_clause 0x1 +; G_GFX1310-NEXT: s_load_b64 s[6:7], s[4:5], 0x34 nv +; G_GFX1310-NEXT: s_load_b128 s[0:3], s[4:5], 0x24 nv +; G_GFX1310-NEXT: s_wait_kmcnt 0x0 +; G_GFX1310-NEXT: v_dual_mov_b32 v0, s6 :: v_dual_mov_b32 v1, s7 +; G_GFX1310-NEXT: buffer_atomic_min_num_f32 v0, v1, s[0:3], null offen +; G_GFX1310-NEXT: s_endpgm main_body: %ret = call float @llvm.amdgcn.raw.ptr.buffer.atomic.fmin.f32(float %data, ptr addrspace(8) %rsrc, i32 %vindex, i32 0, i32 0) ret void @@ -164,6 +186,13 @@ define amdgpu_ps void @raw_ptr_buffer_atomic_min_rtn_f32(ptr addrspace(8) inreg ; GFX1100-NEXT: global_store_b32 v[0:1], v0, off ; GFX1100-NEXT: s_endpgm ; +; GFX1310-LABEL: raw_ptr_buffer_atomic_min_rtn_f32: +; GFX1310: ; %bb.0: ; %main_body +; GFX1310-NEXT: buffer_atomic_min_num_f32 v0, v1, s[0:3], null offen th:TH_ATOMIC_RETURN +; GFX1310-NEXT: s_wait_loadcnt 0x0 +; GFX1310-NEXT: global_store_b32 v[0:1], v0, off +; GFX1310-NEXT: s_endpgm +; ; G_SI-LABEL: raw_ptr_buffer_atomic_min_rtn_f32: ; G_SI: ; %bb.0: ; %main_body ; G_SI-NEXT: buffer_atomic_fmin v0, v1, s[0:3], 0 offen glc @@ -202,6 +231,13 @@ define amdgpu_ps void @raw_ptr_buffer_atomic_min_rtn_f32(ptr addrspace(8) inreg ; G_GFX1100-NEXT: s_waitcnt vmcnt(0) ; G_GFX1100-NEXT: global_store_b32 v[0:1], v0, off ; G_GFX1100-NEXT: s_endpgm +; +; G_GFX1310-LABEL: raw_ptr_buffer_atomic_min_rtn_f32: +; G_GFX1310: ; %bb.0: ; %main_body +; G_GFX1310-NEXT: buffer_atomic_min_num_f32 v0, v1, s[0:3], null offen th:TH_ATOMIC_RETURN +; G_GFX1310-NEXT: s_wait_loadcnt 0x0 +; G_GFX1310-NEXT: global_store_b32 v[0:1], v0, off +; G_GFX1310-NEXT: s_endpgm main_body: %ret = call float @llvm.amdgcn.raw.ptr.buffer.atomic.fmin.f32(float %data, ptr addrspace(8) %rsrc, i32 %vindex, i32 0, i32 0) store float %ret, ptr addrspace(1) poison @@ -273,6 +309,20 @@ define amdgpu_kernel void @raw_ptr_buffer_atomic_min_rtn_f32_off4_slc(ptr addrsp ; GFX1100-NEXT: ds_store_b32 v1, v0 ; GFX1100-NEXT: s_endpgm ; +; GFX1310-LABEL: raw_ptr_buffer_atomic_min_rtn_f32_off4_slc: +; GFX1310: ; %bb.0: ; %main_body +; GFX1310-NEXT: s_clause 0x1 +; GFX1310-NEXT: s_load_b96 s[8:10], s[4:5], 0x34 nv +; GFX1310-NEXT: s_load_b128 s[0:3], s[4:5], 0x24 nv +; GFX1310-NEXT: s_mov_b32 s4, 4 +; GFX1310-NEXT: s_wait_kmcnt 0x0 +; GFX1310-NEXT: v_dual_mov_b32 v0, s8 :: v_dual_mov_b32 v1, s9 +; GFX1310-NEXT: buffer_atomic_min_num_f32 v0, v1, s[0:3], s4 offen th:TH_ATOMIC_NT_RETURN +; GFX1310-NEXT: v_mov_b32_e32 v1, s10 +; GFX1310-NEXT: s_wait_loadcnt 0x0 +; GFX1310-NEXT: ds_store_b32 v1, v0 +; GFX1310-NEXT: s_endpgm +; ; G_SI-LABEL: raw_ptr_buffer_atomic_min_rtn_f32_off4_slc: ; G_SI: ; %bb.0: ; %main_body ; G_SI-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0xd @@ -352,6 +402,20 @@ define amdgpu_kernel void @raw_ptr_buffer_atomic_min_rtn_f32_off4_slc(ptr addrsp ; G_GFX1100-NEXT: s_waitcnt vmcnt(0) ; G_GFX1100-NEXT: ds_store_b32 v1, v0 ; G_GFX1100-NEXT: s_endpgm +; +; G_GFX1310-LABEL: raw_ptr_buffer_atomic_min_rtn_f32_off4_slc: +; G_GFX1310: ; %bb.0: ; %main_body +; G_GFX1310-NEXT: s_clause 0x1 +; G_GFX1310-NEXT: s_load_b96 s[8:10], s[4:5], 0x34 nv +; G_GFX1310-NEXT: s_load_b128 s[0:3], s[4:5], 0x24 nv +; G_GFX1310-NEXT: s_mov_b32 s4, 4 +; G_GFX1310-NEXT: s_wait_kmcnt 0x0 +; G_GFX1310-NEXT: v_dual_mov_b32 v0, s8 :: v_dual_mov_b32 v1, s9 +; G_GFX1310-NEXT: buffer_atomic_min_num_f32 v0, v1, s[0:3], s4 offen th:TH_ATOMIC_NT_RETURN +; G_GFX1310-NEXT: v_mov_b32_e32 v1, s10 +; G_GFX1310-NEXT: s_wait_loadcnt 0x0 +; G_GFX1310-NEXT: ds_store_b32 v1, v0 +; G_GFX1310-NEXT: s_endpgm ; GFX1010-LABEL: raw_ptr_buffer_atomic_min_rtn_f32_off4_slc: main_body: %ret = call float @llvm.amdgcn.raw.ptr.buffer.atomic.fmin.f32(float %data, ptr addrspace(8) %rsrc, i32 %vindex, i32 4, i32 2) @@ -412,6 +476,16 @@ define amdgpu_kernel void @raw_ptr_buffer_atomic_max_noret_f32(ptr addrspace(8) ; GFX1100-NEXT: buffer_atomic_max_f32 v0, v1, s[0:3], 0 offen ; GFX1100-NEXT: s_endpgm ; +; GFX1310-LABEL: raw_ptr_buffer_atomic_max_noret_f32: +; GFX1310: ; %bb.0: ; %main_body +; GFX1310-NEXT: s_clause 0x1 +; GFX1310-NEXT: s_load_b64 s[6:7], s[4:5], 0x34 nv +; GFX1310-NEXT: s_load_b128 s[0:3], s[4:5], 0x24 nv +; GFX1310-NEXT: s_wait_kmcnt 0x0 +; GFX1310-NEXT: v_dual_mov_b32 v0, s6 :: v_dual_mov_b32 v1, s7 +; GFX1310-NEXT: buffer_atomic_max_num_f32 v0, v1, s[0:3], null offen +; GFX1310-NEXT: s_endpgm +; ; G_SI-LABEL: raw_ptr_buffer_atomic_max_noret_f32: ; G_SI: ; %bb.0: ; %main_body ; G_SI-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0xd @@ -463,6 +537,16 @@ define amdgpu_kernel void @raw_ptr_buffer_atomic_max_noret_f32(ptr addrspace(8) ; G_GFX1100-NEXT: v_dual_mov_b32 v0, s6 :: v_dual_mov_b32 v1, s7 ; G_GFX1100-NEXT: buffer_atomic_max_f32 v0, v1, s[0:3], 0 offen ; G_GFX1100-NEXT: s_endpgm +; +; G_GFX1310-LABEL: raw_ptr_buffer_atomic_max_noret_f32: +; G_GFX1310: ; %bb.0: ; %main_body +; G_GFX1310-NEXT: s_clause 0x1 +; G_GFX1310-NEXT: s_load_b64 s[6:7], s[4:5], 0x34 nv +; G_GFX1310-NEXT: s_load_b128 s[0:3], s[4:5], 0x24 nv +; G_GFX1310-NEXT: s_wait_kmcnt 0x0 +; G_GFX1310-NEXT: v_dual_mov_b32 v0, s6 :: v_dual_mov_b32 v1, s7 +; G_GFX1310-NEXT: buffer_atomic_max_num_f32 v0, v1, s[0:3], null offen +; G_GFX1310-NEXT: s_endpgm main_body: %ret = call float @llvm.amdgcn.raw.ptr.buffer.atomic.fmax.f32(float %data, ptr addrspace(8) %rsrc, i32 %vindex, i32 0, i32 0) ret void @@ -508,6 +592,13 @@ define amdgpu_ps void @raw_ptr_buffer_atomic_max_rtn_f32(ptr addrspace(8) inreg ; GFX1100-NEXT: global_store_b32 v[0:1], v0, off ; GFX1100-NEXT: s_endpgm ; +; GFX1310-LABEL: raw_ptr_buffer_atomic_max_rtn_f32: +; GFX1310: ; %bb.0: ; %main_body +; GFX1310-NEXT: buffer_atomic_max_num_f32 v0, v1, s[0:3], null offen th:TH_ATOMIC_RETURN +; GFX1310-NEXT: s_wait_loadcnt 0x0 +; GFX1310-NEXT: global_store_b32 v[0:1], v0, off +; GFX1310-NEXT: s_endpgm +; ; G_SI-LABEL: raw_ptr_buffer_atomic_max_rtn_f32: ; G_SI: ; %bb.0: ; %main_body ; G_SI-NEXT: buffer_atomic_fmax v0, v1, s[0:3], 0 offen glc @@ -546,6 +637,13 @@ define amdgpu_ps void @raw_ptr_buffer_atomic_max_rtn_f32(ptr addrspace(8) inreg ; G_GFX1100-NEXT: s_waitcnt vmcnt(0) ; G_GFX1100-NEXT: global_store_b32 v[0:1], v0, off ; G_GFX1100-NEXT: s_endpgm +; +; G_GFX1310-LABEL: raw_ptr_buffer_atomic_max_rtn_f32: +; G_GFX1310: ; %bb.0: ; %main_body +; G_GFX1310-NEXT: buffer_atomic_max_num_f32 v0, v1, s[0:3], null offen th:TH_ATOMIC_RETURN +; G_GFX1310-NEXT: s_wait_loadcnt 0x0 +; G_GFX1310-NEXT: global_store_b32 v[0:1], v0, off +; G_GFX1310-NEXT: s_endpgm main_body: %ret = call float @llvm.amdgcn.raw.ptr.buffer.atomic.fmax.f32(float %data, ptr addrspace(8) %rsrc, i32 %vindex, i32 0, i32 0) store float %ret, ptr addrspace(1) poison @@ -618,6 +716,18 @@ define amdgpu_kernel void @raw_ptr_buffer_atomic_max_rtn_f32_off4_slc(ptr addrsp ; GFX1100-NEXT: global_store_b32 v1, v0, s[6:7] ; GFX1100-NEXT: s_endpgm ; +; GFX1310-LABEL: raw_ptr_buffer_atomic_max_rtn_f32_off4_slc: +; GFX1310: ; %bb.0: ; %main_body +; GFX1310-NEXT: s_load_b256 s[0:7], s[4:5], 0x24 nv +; GFX1310-NEXT: s_wait_kmcnt 0x0 +; GFX1310-NEXT: v_dual_mov_b32 v0, s4 :: v_dual_mov_b32 v1, s5 +; GFX1310-NEXT: s_mov_b32 s4, 4 +; GFX1310-NEXT: buffer_atomic_max_num_f32 v0, v1, s[0:3], s4 offen th:TH_ATOMIC_NT_RETURN +; GFX1310-NEXT: v_mov_b32_e32 v1, 0 +; GFX1310-NEXT: s_wait_loadcnt 0x0 +; GFX1310-NEXT: global_store_b32 v1, v0, s[6:7] +; GFX1310-NEXT: s_endpgm +; ; G_SI-LABEL: raw_ptr_buffer_atomic_max_rtn_f32_off4_slc: ; G_SI: ; %bb.0: ; %main_body ; G_SI-NEXT: s_load_dwordx8 s[0:7], s[4:5], 0x9 @@ -680,6 +790,18 @@ define amdgpu_kernel void @raw_ptr_buffer_atomic_max_rtn_f32_off4_slc(ptr addrsp ; G_GFX1100-NEXT: s_waitcnt vmcnt(0) ; G_GFX1100-NEXT: global_store_b32 v1, v0, s[6:7] ; G_GFX1100-NEXT: s_endpgm +; +; G_GFX1310-LABEL: raw_ptr_buffer_atomic_max_rtn_f32_off4_slc: +; G_GFX1310: ; %bb.0: ; %main_body +; G_GFX1310-NEXT: s_load_b256 s[0:7], s[4:5], 0x24 nv +; G_GFX1310-NEXT: s_wait_kmcnt 0x0 +; G_GFX1310-NEXT: v_dual_mov_b32 v0, s4 :: v_dual_mov_b32 v1, s5 +; G_GFX1310-NEXT: s_mov_b32 s4, 4 +; G_GFX1310-NEXT: buffer_atomic_max_num_f32 v0, v1, s[0:3], s4 offen th:TH_ATOMIC_NT_RETURN +; G_GFX1310-NEXT: v_mov_b32_e32 v1, 0 +; G_GFX1310-NEXT: s_wait_loadcnt 0x0 +; G_GFX1310-NEXT: global_store_b32 v1, v0, s[6:7] +; G_GFX1310-NEXT: s_endpgm main_body: %ret = call float @llvm.amdgcn.raw.ptr.buffer.atomic.fmax.f32(float %data, ptr addrspace(8) %rsrc, i32 %vindex, i32 4, i32 2) store float %ret, ptr addrspace(1) %out, align 8 diff --git a/llvm/test/CodeGen/AMDGPU/fp64-min-max-buffer-ptr-atomics.ll b/llvm/test/CodeGen/AMDGPU/fp64-min-max-buffer-ptr-atomics.ll index a6791f7ddf9ea..e9487127b23f3 100644 --- a/llvm/test/CodeGen/AMDGPU/fp64-min-max-buffer-ptr-atomics.ll +++ b/llvm/test/CodeGen/AMDGPU/fp64-min-max-buffer-ptr-atomics.ll @@ -3,11 +3,13 @@ ; RUN: llc < %s -mtriple=amdgpu7.01 | FileCheck %s -check-prefix=GFX7 ; RUN: llc < %s -mtriple=amdgpu10.10 | FileCheck %s -check-prefix=GFX10 ; RUN: llc < %s -mtriple=amdgpu10.30 | FileCheck %s -check-prefix=GFX1030 +; RUN: llc < %s -mtriple=amdgpu13.10 | FileCheck %s -check-prefix=GFX1310 ; RUN: llc < %s -global-isel -mtriple=amdgpu6.01 | FileCheck %s -check-prefix=G_SI ; RUN: llc < %s -global-isel -mtriple=amdgpu7.01 | FileCheck %s -check-prefix=G_GFX7 ; RUN: llc < %s -global-isel -mtriple=amdgpu10.10 | FileCheck %s -check-prefix=G_GFX10 ; RUN: llc < %s -global-isel -mtriple=amdgpu10.30 | FileCheck %s -check-prefix=G_GFX1030 +; RUN: llc < %s -global-isel -mtriple=amdgpu13.10 | FileCheck %s -check-prefix=G_GFX1310 declare double @llvm.amdgcn.raw.ptr.buffer.atomic.fmin.f64(double, ptr addrspace(8), i32, i32, i32 immarg) declare double @llvm.amdgcn.raw.ptr.buffer.atomic.fmax.f64(double, ptr addrspace(8), i32, i32, i32 immarg) @@ -64,6 +66,17 @@ define amdgpu_kernel void @raw_ptr_buffer_atomic_min_noret_f64(ptr addrspace(8) ; GFX1030-NEXT: buffer_atomic_fmin_x2 v[0:1], v2, s[0:3], 0 offen ; GFX1030-NEXT: s_endpgm ; +; GFX1310-LABEL: raw_ptr_buffer_atomic_min_noret_f64: +; GFX1310: ; %bb.0: ; %main_body +; GFX1310-NEXT: s_clause 0x1 +; GFX1310-NEXT: s_load_b96 s[8:10], s[4:5], 0x34 nv +; GFX1310-NEXT: s_load_b128 s[0:3], s[4:5], 0x24 nv +; GFX1310-NEXT: s_wait_kmcnt 0x0 +; GFX1310-NEXT: v_dual_mov_b32 v0, s8 :: v_dual_mov_b32 v1, s9 +; GFX1310-NEXT: v_mov_b32_e32 v2, s10 +; GFX1310-NEXT: buffer_atomic_min_num_f64 v[0:1], v2, s[0:3], null offen +; GFX1310-NEXT: s_endpgm +; ; G_SI-LABEL: raw_ptr_buffer_atomic_min_noret_f64: ; G_SI: ; %bb.0: ; %main_body ; G_SI-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0xd @@ -113,6 +126,17 @@ define amdgpu_kernel void @raw_ptr_buffer_atomic_min_noret_f64(ptr addrspace(8) ; G_GFX1030-NEXT: v_mov_b32_e32 v2, s8 ; G_GFX1030-NEXT: buffer_atomic_fmin_x2 v[0:1], v2, s[0:3], 0 offen ; G_GFX1030-NEXT: s_endpgm +; +; G_GFX1310-LABEL: raw_ptr_buffer_atomic_min_noret_f64: +; G_GFX1310: ; %bb.0: ; %main_body +; G_GFX1310-NEXT: s_clause 0x1 +; G_GFX1310-NEXT: s_load_b96 s[8:10], s[4:5], 0x34 nv +; G_GFX1310-NEXT: s_load_b128 s[0:3], s[4:5], 0x24 nv +; G_GFX1310-NEXT: s_wait_kmcnt 0x0 +; G_GFX1310-NEXT: v_dual_mov_b32 v0, s8 :: v_dual_mov_b32 v1, s9 +; G_GFX1310-NEXT: v_mov_b32_e32 v2, s10 +; G_GFX1310-NEXT: buffer_atomic_min_num_f64 v[0:1], v2, s[0:3], null offen +; G_GFX1310-NEXT: s_endpgm main_body: %ret = call double @llvm.amdgcn.raw.ptr.buffer.atomic.fmin.f64(double %data, ptr addrspace(8) %rsrc, i32 %vindex, i32 0, i32 0) ret void @@ -149,6 +173,13 @@ define amdgpu_ps void @raw_ptr_buffer_atomic_min_rtn_f64(ptr addrspace(8) inreg ; GFX1030-NEXT: ds_write_b64 v0, v[0:1] ; GFX1030-NEXT: s_endpgm ; +; GFX1310-LABEL: raw_ptr_buffer_atomic_min_rtn_f64: +; GFX1310: ; %bb.0: ; %main_body +; GFX1310-NEXT: buffer_atomic_min_num_f64 v[0:1], v2, s[0:3], null offen th:TH_ATOMIC_RETURN +; GFX1310-NEXT: s_wait_loadcnt 0x0 +; GFX1310-NEXT: ds_store_b64 v0, v[0:1] +; GFX1310-NEXT: s_endpgm +; ; G_SI-LABEL: raw_ptr_buffer_atomic_min_rtn_f64: ; G_SI: ; %bb.0: ; %main_body ; G_SI-NEXT: buffer_atomic_fmin_x2 v[0:1], v2, s[0:3], 0 offen glc @@ -178,6 +209,13 @@ define amdgpu_ps void @raw_ptr_buffer_atomic_min_rtn_f64(ptr addrspace(8) inreg ; G_GFX1030-NEXT: s_waitcnt vmcnt(0) ; G_GFX1030-NEXT: ds_write_b64 v0, v[0:1] ; G_GFX1030-NEXT: s_endpgm +; +; G_GFX1310-LABEL: raw_ptr_buffer_atomic_min_rtn_f64: +; G_GFX1310: ; %bb.0: ; %main_body +; G_GFX1310-NEXT: buffer_atomic_min_num_f64 v[0:1], v2, s[0:3], null offen th:TH_ATOMIC_RETURN +; G_GFX1310-NEXT: s_wait_loadcnt 0x0 +; G_GFX1310-NEXT: ds_store_b64 v0, v[0:1] +; G_GFX1310-NEXT: s_endpgm main_body: %ret = call double @llvm.amdgcn.raw.ptr.buffer.atomic.fmin.f64(double %data, ptr addrspace(8) %rsrc, i32 %vindex, i32 0, i32 0) store double %ret, ptr addrspace(3) poison @@ -215,6 +253,14 @@ define amdgpu_ps void @raw_ptr_buffer_atomic_min_rtn_f64_off4_slc(ptr addrspace( ; GFX1030-NEXT: ds_write_b64 v3, v[0:1] ; GFX1030-NEXT: s_endpgm ; +; GFX1310-LABEL: raw_ptr_buffer_atomic_min_rtn_f64_off4_slc: +; GFX1310: ; %bb.0: ; %main_body +; GFX1310-NEXT: s_mov_b32 s4, 4 +; GFX1310-NEXT: buffer_atomic_min_num_f64 v[0:1], v2, s[0:3], s4 offen th:TH_ATOMIC_NT_RETURN +; GFX1310-NEXT: s_wait_loadcnt 0x0 +; GFX1310-NEXT: ds_store_b64 v3, v[0:1] +; GFX1310-NEXT: s_endpgm +; ; G_SI-LABEL: raw_ptr_buffer_atomic_min_rtn_f64_off4_slc: ; G_SI: ; %bb.0: ; %main_body ; G_SI-NEXT: buffer_atomic_fmin_x2 v[0:1], v2, s[0:3], 4 offen glc slc @@ -244,6 +290,14 @@ define amdgpu_ps void @raw_ptr_buffer_atomic_min_rtn_f64_off4_slc(ptr addrspace( ; G_GFX1030-NEXT: s_waitcnt vmcnt(0) ; G_GFX1030-NEXT: ds_write_b64 v3, v[0:1] ; G_GFX1030-NEXT: s_endpgm +; +; G_GFX1310-LABEL: raw_ptr_buffer_atomic_min_rtn_f64_off4_slc: +; G_GFX1310: ; %bb.0: ; %main_body +; G_GFX1310-NEXT: s_mov_b32 s4, 4 +; G_GFX1310-NEXT: buffer_atomic_min_num_f64 v[0:1], v2, s[0:3], s4 offen th:TH_ATOMIC_NT_RETURN +; G_GFX1310-NEXT: s_wait_loadcnt 0x0 +; G_GFX1310-NEXT: ds_store_b64 v3, v[0:1] +; G_GFX1310-NEXT: s_endpgm main_body: %ret = call double @llvm.amdgcn.raw.ptr.buffer.atomic.fmin.f64(double %data, ptr addrspace(8) %rsrc, i32 %vindex, i32 4, i32 2) store double %ret, ptr addrspace(3) %out, align 8 @@ -301,6 +355,17 @@ define amdgpu_kernel void @raw_ptr_buffer_atomic_max_noret_f64(ptr addrspace(8) ; GFX1030-NEXT: buffer_atomic_fmax_x2 v[0:1], v2, s[0:3], 0 offen ; GFX1030-NEXT: s_endpgm ; +; GFX1310-LABEL: raw_ptr_buffer_atomic_max_noret_f64: +; GFX1310: ; %bb.0: ; %main_body +; GFX1310-NEXT: s_clause 0x1 +; GFX1310-NEXT: s_load_b96 s[8:10], s[4:5], 0x34 nv +; GFX1310-NEXT: s_load_b128 s[0:3], s[4:5], 0x24 nv +; GFX1310-NEXT: s_wait_kmcnt 0x0 +; GFX1310-NEXT: v_dual_mov_b32 v0, s8 :: v_dual_mov_b32 v1, s9 +; GFX1310-NEXT: v_mov_b32_e32 v2, s10 +; GFX1310-NEXT: buffer_atomic_max_num_f64 v[0:1], v2, s[0:3], null offen +; GFX1310-NEXT: s_endpgm +; ; G_SI-LABEL: raw_ptr_buffer_atomic_max_noret_f64: ; G_SI: ; %bb.0: ; %main_body ; G_SI-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0xd @@ -350,6 +415,17 @@ define amdgpu_kernel void @raw_ptr_buffer_atomic_max_noret_f64(ptr addrspace(8) ; G_GFX1030-NEXT: v_mov_b32_e32 v2, s8 ; G_GFX1030-NEXT: buffer_atomic_fmax_x2 v[0:1], v2, s[0:3], 0 offen ; G_GFX1030-NEXT: s_endpgm +; +; G_GFX1310-LABEL: raw_ptr_buffer_atomic_max_noret_f64: +; G_GFX1310: ; %bb.0: ; %main_body +; G_GFX1310-NEXT: s_clause 0x1 +; G_GFX1310-NEXT: s_load_b96 s[8:10], s[4:5], 0x34 nv +; G_GFX1310-NEXT: s_load_b128 s[0:3], s[4:5], 0x24 nv +; G_GFX1310-NEXT: s_wait_kmcnt 0x0 +; G_GFX1310-NEXT: v_dual_mov_b32 v0, s8 :: v_dual_mov_b32 v1, s9 +; G_GFX1310-NEXT: v_mov_b32_e32 v2, s10 +; G_GFX1310-NEXT: buffer_atomic_max_num_f64 v[0:1], v2, s[0:3], null offen +; G_GFX1310-NEXT: s_endpgm main_body: %ret = call double @llvm.amdgcn.raw.ptr.buffer.atomic.fmax.f64(double %data, ptr addrspace(8) %rsrc, i32 %vindex, i32 0, i32 0) ret void @@ -386,6 +462,13 @@ define amdgpu_ps void @raw_ptr_buffer_atomic_max_rtn_f64(ptr addrspace(8) inreg ; GFX1030-NEXT: ds_write_b64 v0, v[0:1] ; GFX1030-NEXT: s_endpgm ; +; GFX1310-LABEL: raw_ptr_buffer_atomic_max_rtn_f64: +; GFX1310: ; %bb.0: ; %main_body +; GFX1310-NEXT: buffer_atomic_max_num_f64 v[0:1], v2, s[0:3], null offen th:TH_ATOMIC_RETURN +; GFX1310-NEXT: s_wait_loadcnt 0x0 +; GFX1310-NEXT: ds_store_b64 v0, v[0:1] +; GFX1310-NEXT: s_endpgm +; ; G_SI-LABEL: raw_ptr_buffer_atomic_max_rtn_f64: ; G_SI: ; %bb.0: ; %main_body ; G_SI-NEXT: buffer_atomic_fmax_x2 v[0:1], v2, s[0:3], 0 offen glc @@ -415,6 +498,13 @@ define amdgpu_ps void @raw_ptr_buffer_atomic_max_rtn_f64(ptr addrspace(8) inreg ; G_GFX1030-NEXT: s_waitcnt vmcnt(0) ; G_GFX1030-NEXT: ds_write_b64 v0, v[0:1] ; G_GFX1030-NEXT: s_endpgm +; +; G_GFX1310-LABEL: raw_ptr_buffer_atomic_max_rtn_f64: +; G_GFX1310: ; %bb.0: ; %main_body +; G_GFX1310-NEXT: buffer_atomic_max_num_f64 v[0:1], v2, s[0:3], null offen th:TH_ATOMIC_RETURN +; G_GFX1310-NEXT: s_wait_loadcnt 0x0 +; G_GFX1310-NEXT: ds_store_b64 v0, v[0:1] +; G_GFX1310-NEXT: s_endpgm main_body: %ret = call double @llvm.amdgcn.raw.ptr.buffer.atomic.fmax.f64(double %data, ptr addrspace(8) %rsrc, i32 %vindex, i32 0, i32 0) store double %ret, ptr addrspace(3) poison @@ -476,6 +566,19 @@ define amdgpu_kernel void @raw_ptr_buffer_atomic_max_rtn_f64_off4_slc(ptr addrsp ; GFX1030-NEXT: ds_write_b64 v2, v[0:1] ; GFX1030-NEXT: s_endpgm ; +; GFX1310-LABEL: raw_ptr_buffer_atomic_max_rtn_f64_off4_slc: +; GFX1310: ; %bb.0: ; %main_body +; GFX1310-NEXT: s_load_b256 s[0:7], s[4:5], 0x24 nv +; GFX1310-NEXT: s_wait_kmcnt 0x0 +; GFX1310-NEXT: v_dual_mov_b32 v0, s4 :: v_dual_mov_b32 v1, s5 +; GFX1310-NEXT: v_mov_b32_e32 v2, s6 +; GFX1310-NEXT: s_mov_b32 s4, 4 +; GFX1310-NEXT: buffer_atomic_max_num_f64 v[0:1], v2, s[0:3], s4 offen th:TH_ATOMIC_NT_RETURN +; GFX1310-NEXT: v_mov_b32_e32 v2, s7 +; GFX1310-NEXT: s_wait_loadcnt 0x0 +; GFX1310-NEXT: ds_store_b64 v2, v[0:1] +; GFX1310-NEXT: s_endpgm +; ; G_SI-LABEL: raw_ptr_buffer_atomic_max_rtn_f64_off4_slc: ; G_SI: ; %bb.0: ; %main_body ; G_SI-NEXT: s_load_dwordx8 s[0:7], s[4:5], 0x9 @@ -529,6 +632,19 @@ define amdgpu_kernel void @raw_ptr_buffer_atomic_max_rtn_f64_off4_slc(ptr addrsp ; G_GFX1030-NEXT: s_waitcnt vmcnt(0) ; G_GFX1030-NEXT: ds_write_b64 v2, v[0:1] ; G_GFX1030-NEXT: s_endpgm +; +; G_GFX1310-LABEL: raw_ptr_buffer_atomic_max_rtn_f64_off4_slc: +; G_GFX1310: ; %bb.0: ; %main_body +; G_GFX1310-NEXT: s_load_b256 s[0:7], s[4:5], 0x24 nv +; G_GFX1310-NEXT: s_wait_kmcnt 0x0 +; G_GFX1310-NEXT: v_dual_mov_b32 v0, s4 :: v_dual_mov_b32 v1, s5 +; G_GFX1310-NEXT: v_mov_b32_e32 v2, s6 +; G_GFX1310-NEXT: s_mov_b32 s4, 4 +; G_GFX1310-NEXT: buffer_atomic_max_num_f64 v[0:1], v2, s[0:3], s4 offen th:TH_ATOMIC_NT_RETURN +; G_GFX1310-NEXT: v_mov_b32_e32 v2, s7 +; G_GFX1310-NEXT: s_wait_loadcnt 0x0 +; G_GFX1310-NEXT: ds_store_b64 v2, v[0:1] +; G_GFX1310-NEXT: s_endpgm main_body: %ret = call double @llvm.amdgcn.raw.ptr.buffer.atomic.fmax.f64(double %data, ptr addrspace(8) %rsrc, i32 %vindex, i32 4, i32 2) store double %ret, ptr addrspace(3) %out, align 8 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
