https://github.com/arsenm created
https://github.com/llvm/llvm-project/pull/213119
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)
>From a2b26bb6c0e231bd3ba05994f94429052192f569 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 | 38 ++++++
.../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, 279 insertions(+)
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl
index a6e0b15921860..715f3af93bd15 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx13.cl
@@ -46,3 +46,41 @@ 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(
+// 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 double, ptr addrspace(5) [[D_ADDR]],
align 8
+// 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 double
@llvm.amdgcn.raw.ptr.buffer.atomic.fmin.f64(double [[TMP6]], ptr addrspace(8)
[[TMP7]], i32 [[TMP8]], i32 [[TMP9]], i32 0)
+// CHECK-NEXT: [[TMP11:%.*]] = load ptr addrspace(1), ptr addrspace(5)
[[DOUT_ADDR]], align 8
+// CHECK-NEXT: store double [[TMP10]], ptr addrspace(1) [[TMP11]], align 8
+// CHECK-NEXT: ret void
+//
+void test_raw_ptr_buffer_atomic_fmin(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);
+ *dout = __builtin_amdgcn_raw_ptr_buffer_atomic_fmin_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 6e16bc8dc9507..a82a20bf3a41c 100644
--- a/llvm/lib/TargetParser/AMDGPUTargetParser.cpp
+++ b/llvm/lib/TargetParser/AMDGPUTargetParser.cpp
@@ -516,6 +516,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
