[llvm] [AMDGPU] Guard media instruction patterns with target feature (PR #217995)
Shilei Tian via llvm-commits
llvm-commits at lists.llvm.org
Fri Aug 21 11:53:21 PDT 2026
https://github.com/shiltian updated https://github.com/llvm/llvm-project/pull/217995
>From dbee7c9a000dfd80b989ba86cda316324f839a4d Mon Sep 17 00:00:00 2001
From: Shilei Tian <i at tianshilei.me>
Date: Fri, 21 Aug 2026 14:07:25 -0400
Subject: [PATCH] [AMDGPU] Guard media instruction patterns with target feature
As well as their corresponding intrinsics.
---
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 14 +-
llvm/lib/Target/AMDGPU/SIInstructions.td | 28 +-
.../AMDGPU/llvm.amdgcn.intrinsics.sad.err.ll | 44 +
llvm/test/CodeGen/AMDGPU/sad.ll | 1965 ++++++++++++++---
4 files changed, 1677 insertions(+), 374 deletions(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/llvm.amdgcn.intrinsics.sad.err.ll
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index b321b9155a199..191501b7862cb 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -2471,14 +2471,11 @@ def int_amdgcn_lerp :
ClangBuiltin<"__builtin_amdgcn_lerp">,
PureIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty]>;
+let TargetFeatures = "sad-insts" in {
def int_amdgcn_sad_u8 :
ClangBuiltin<"__builtin_amdgcn_sad_u8">,
PureIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty]>;
-def int_amdgcn_msad_u8 :
- ClangBuiltin<"__builtin_amdgcn_msad_u8">,
- PureIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty]>;
-
def int_amdgcn_sad_hi_u8 :
ClangBuiltin<"__builtin_amdgcn_sad_hi_u8">,
PureIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty]>;
@@ -2486,15 +2483,24 @@ def int_amdgcn_sad_hi_u8 :
def int_amdgcn_sad_u16 :
ClangBuiltin<"__builtin_amdgcn_sad_u16">,
PureIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty]>;
+} // End TargetFeatures = "sad-insts"
+
+let TargetFeatures = "msad-insts" in
+def int_amdgcn_msad_u8 :
+ ClangBuiltin<"__builtin_amdgcn_msad_u8">,
+ PureIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty]>;
+let TargetFeatures = "qsad-insts" in
def int_amdgcn_qsad_pk_u16_u8 :
ClangBuiltin<"__builtin_amdgcn_qsad_pk_u16_u8">,
PureIntrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty]>;
+let TargetFeatures = "mqsad-pk-insts" in
def int_amdgcn_mqsad_pk_u16_u8 :
ClangBuiltin<"__builtin_amdgcn_mqsad_pk_u16_u8">,
PureIntrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty]>;
+let TargetFeatures = "mqsad-insts" in
def int_amdgcn_mqsad_u32_u8 :
ClangBuiltin<"__builtin_amdgcn_mqsad_u32_u8">,
PureIntrinsic<[llvm_v4i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_v4i32_ty]>;
diff --git a/llvm/lib/Target/AMDGPU/SIInstructions.td b/llvm/lib/Target/AMDGPU/SIInstructions.td
index c2fe750371528..441098168f450 100644
--- a/llvm/lib/Target/AMDGPU/SIInstructions.td
+++ b/llvm/lib/Target/AMDGPU/SIInstructions.td
@@ -3100,20 +3100,22 @@ defm : SI_INDIRECT_Pattern <v32i32, i32, "V32">;
// SAD Patterns
//===----------------------------------------------------------------------===//
-def : GCNPat <
- (add (sub_oneuse (umax i32:$src0, i32:$src1),
- (umin i32:$src0, i32:$src1)),
- i32:$src2),
- (V_SAD_U32_e64 $src0, $src1, $src2, (i1 0))
->;
+let SubtargetPredicate = HasSadInsts in {
+ def : GCNPat <
+ (add (sub_oneuse (umax i32:$src0, i32:$src1),
+ (umin i32:$src0, i32:$src1)),
+ i32:$src2),
+ (V_SAD_U32_e64 $src0, $src1, $src2, (i1 0))
+ >;
-def : GCNPat <
- (add (select_oneuse (i1 (setugt i32:$src0, i32:$src1)),
- (sub i32:$src0, i32:$src1),
- (sub i32:$src1, i32:$src0)),
- i32:$src2),
- (V_SAD_U32_e64 $src0, $src1, $src2, (i1 0))
->;
+ def : GCNPat <
+ (add (select_oneuse (i1 (setugt i32:$src0, i32:$src1)),
+ (sub i32:$src0, i32:$src1),
+ (sub i32:$src1, i32:$src0)),
+ i32:$src2),
+ (V_SAD_U32_e64 $src0, $src1, $src2, (i1 0))
+ >;
+} // End SubtargetPredicate = HasSadInsts
//===----------------------------------------------------------------------===//
// Conversion Patterns
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.intrinsics.sad.err.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.intrinsics.sad.err.ll
new file mode 100644
index 0000000000000..3135b7953f60d
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.intrinsics.sad.err.ll
@@ -0,0 +1,44 @@
+; RUN: not llc -global-isel=0 -mtriple=amdgpu12.5-amd-amdhsa -filetype=null < %s 2>&1 | FileCheck %s
+; RUN: not llc -global-isel=1 -global-isel-abort=0 -mtriple=amdgpu12.5-amd-amdhsa -filetype=null < %s 2>&1 | FileCheck %s
+
+; CHECK: error: {{.*}} llvm.amdgcn.sad.u8 requires target feature 'sad-insts'
+define i32 @test_sad_u8(i32 %a, i32 %b, i32 %c) {
+ %result = call i32 @llvm.amdgcn.sad.u8(i32 %a, i32 %b, i32 %c)
+ ret i32 %result
+}
+
+; CHECK: error: {{.*}} llvm.amdgcn.sad.hi.u8 requires target feature 'sad-insts'
+define i32 @test_sad_hi_u8(i32 %a, i32 %b, i32 %c) {
+ %result = call i32 @llvm.amdgcn.sad.hi.u8(i32 %a, i32 %b, i32 %c)
+ ret i32 %result
+}
+
+; CHECK: error: {{.*}} llvm.amdgcn.sad.u16 requires target feature 'sad-insts'
+define i32 @test_sad_u16(i32 %a, i32 %b, i32 %c) {
+ %result = call i32 @llvm.amdgcn.sad.u16(i32 %a, i32 %b, i32 %c)
+ ret i32 %result
+}
+
+; CHECK: error: {{.*}} llvm.amdgcn.msad.u8 requires target feature 'msad-insts'
+define i32 @test_msad_u8(i32 %a, i32 %b, i32 %c) {
+ %result = call i32 @llvm.amdgcn.msad.u8(i32 %a, i32 %b, i32 %c)
+ ret i32 %result
+}
+
+; CHECK: error: {{.*}} llvm.amdgcn.qsad.pk.u16.u8 requires target feature 'qsad-insts'
+define i64 @test_qsad_pk_u16_u8(i64 %a, i32 %b, i64 %c) {
+ %result = call i64 @llvm.amdgcn.qsad.pk.u16.u8(i64 %a, i32 %b, i64 %c)
+ ret i64 %result
+}
+
+; CHECK: error: {{.*}} llvm.amdgcn.mqsad.pk.u16.u8 requires target feature 'mqsad-pk-insts'
+define i64 @test_mqsad_pk_u16_u8(i64 %a, i32 %b, i64 %c) {
+ %result = call i64 @llvm.amdgcn.mqsad.pk.u16.u8(i64 %a, i32 %b, i64 %c)
+ ret i64 %result
+}
+
+; CHECK: error: {{.*}} llvm.amdgcn.mqsad.u32.u8 requires target feature 'mqsad-insts'
+define <4 x i32> @test_mqsad_u32_u8(i64 %a, i32 %b, <4 x i32> %c) {
+ %result = call <4 x i32> @llvm.amdgcn.mqsad.u32.u8(i64 %a, i32 %b, <4 x i32> %c)
+ ret <4 x i32> %result
+}
diff --git a/llvm/test/CodeGen/AMDGPU/sad.ll b/llvm/test/CodeGen/AMDGPU/sad.ll
index b1fd6ce3d7c2c..3b45835ca10d1 100644
--- a/llvm/test/CodeGen/AMDGPU/sad.ll
+++ b/llvm/test/CodeGen/AMDGPU/sad.ll
@@ -1,22 +1,78 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 4
-; RUN: llc -mtriple=amdgpu7.00-amd-amdhsa -earlycse-debug-hash < %s | FileCheck -check-prefix=GCN %s
+; RUN: llc -global-isel=0 -mtriple=amdgpu7.00 < %s | FileCheck -check-prefix=GFX7-SDAG %s
+; RUN: llc -global-isel=1 -mtriple=amdgpu7.00 < %s | FileCheck -check-prefix=GFX7-GISEL %s
+; RUN: llc -global-isel=0 -mtriple=amdgpu12.5 < %s | FileCheck -check-prefix=GFX12-5-SDAG %s
+; RUN: llc -global-isel=1 -mtriple=amdgpu12.5 < %s | FileCheck -check-prefix=GFX12-5-GISEL %s
define amdgpu_kernel void @v_sad_u32_pat1(ptr addrspace(1) %out, i32 %a, i32 %b, i32 %c) {
-; GCN-LABEL: v_sad_u32_pat1:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_load_dwordx4 s[0:3], s[8:9], 0x2
-; GCN-NEXT: s_load_dwordx2 s[4:5], s[8:9], 0x0
-; GCN-NEXT: s_add_i32 s12, s12, s17
-; GCN-NEXT: s_mov_b32 flat_scratch_lo, s13
-; GCN-NEXT: s_lshr_b32 flat_scratch_hi, s12, 8
-; GCN-NEXT: s_waitcnt lgkmcnt(0)
-; GCN-NEXT: v_mov_b32_e32 v0, s1
-; GCN-NEXT: v_mov_b32_e32 v1, s2
-; GCN-NEXT: v_sad_u32 v2, s0, v0, v1
-; GCN-NEXT: v_mov_b32_e32 v0, s4
-; GCN-NEXT: v_mov_b32_e32 v1, s5
-; GCN-NEXT: flat_store_dword v[0:1], v2
-; GCN-NEXT: s_endpgm
+; GFX7-SDAG-LABEL: v_sad_u32_pat1:
+; GFX7-SDAG: ; %bb.0:
+; GFX7-SDAG-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0xb
+; GFX7-SDAG-NEXT: s_load_dwordx2 s[4:5], s[4:5], 0x9
+; GFX7-SDAG-NEXT: s_mov_b32 s7, 0xf000
+; GFX7-SDAG-NEXT: s_mov_b32 s6, -1
+; GFX7-SDAG-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v0, s1
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v1, s2
+; GFX7-SDAG-NEXT: v_sad_u32 v0, s0, v0, v1
+; GFX7-SDAG-NEXT: buffer_store_dword v0, off, s[4:7], 0
+; GFX7-SDAG-NEXT: s_endpgm
+;
+; GFX7-GISEL-LABEL: v_sad_u32_pat1:
+; GFX7-GISEL: ; %bb.0:
+; GFX7-GISEL-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x9
+; GFX7-GISEL-NEXT: s_load_dword s4, s[4:5], 0xd
+; GFX7-GISEL-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-GISEL-NEXT: s_max_u32 s5, s2, s3
+; GFX7-GISEL-NEXT: s_min_u32 s3, s2, s3
+; GFX7-GISEL-NEXT: s_sub_i32 s3, s5, s3
+; GFX7-GISEL-NEXT: s_add_i32 s3, s3, s4
+; GFX7-GISEL-NEXT: s_mov_b32 s2, -1
+; GFX7-GISEL-NEXT: v_mov_b32_e32 v0, s3
+; GFX7-GISEL-NEXT: s_mov_b32 s3, 0xf000
+; GFX7-GISEL-NEXT: buffer_store_dword v0, off, s[0:3], 0
+; GFX7-GISEL-NEXT: s_endpgm
+;
+; GFX12-5-SDAG-LABEL: v_sad_u32_pat1:
+; GFX12-5-SDAG: ; %bb.0:
+; GFX12-5-SDAG-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-SDAG-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-SDAG-NEXT: v_nop
+; GFX12-5-SDAG-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-SDAG-NEXT: s_clause 0x1
+; GFX12-5-SDAG-NEXT: s_load_b96 s[0:2], s[4:5], 0x2c nv
+; GFX12-5-SDAG-NEXT: s_load_b64 s[6:7], s[4:5], 0x24 nv
+; GFX12-5-SDAG-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-SDAG-NEXT: s_min_u32 s3, s0, s1
+; GFX12-5-SDAG-NEXT: s_max_u32 s0, s0, s1
+; GFX12-5-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
+; GFX12-5-SDAG-NEXT: s_sub_co_i32 s0, s0, s3
+; GFX12-5-SDAG-NEXT: s_add_co_i32 s0, s0, s2
+; GFX12-5-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
+; GFX12-5-SDAG-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s0
+; GFX12-5-SDAG-NEXT: global_store_b32 v0, v1, s[6:7]
+; GFX12-5-SDAG-NEXT: s_endpgm
+;
+; GFX12-5-GISEL-LABEL: v_sad_u32_pat1:
+; GFX12-5-GISEL: ; %bb.0:
+; GFX12-5-GISEL-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-GISEL-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-GISEL-NEXT: v_nop
+; GFX12-5-GISEL-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-GISEL-NEXT: s_clause 0x1
+; GFX12-5-GISEL-NEXT: s_load_b96 s[0:2], s[4:5], 0x2c nv
+; GFX12-5-GISEL-NEXT: s_load_b64 s[6:7], s[4:5], 0x24 nv
+; GFX12-5-GISEL-NEXT: v_mov_b32_e32 v1, 0
+; GFX12-5-GISEL-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-GISEL-NEXT: s_max_u32 s3, s0, s1
+; GFX12-5-GISEL-NEXT: s_min_u32 s0, s0, s1
+; GFX12-5-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
+; GFX12-5-GISEL-NEXT: s_sub_co_i32 s0, s3, s0
+; GFX12-5-GISEL-NEXT: s_add_co_i32 s0, s0, s2
+; GFX12-5-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
+; GFX12-5-GISEL-NEXT: v_mov_b32_e32 v0, s0
+; GFX12-5-GISEL-NEXT: global_store_b32 v1, v0, s[6:7]
+; GFX12-5-GISEL-NEXT: s_endpgm
%icmp0 = icmp ugt i32 %a, %b
%t0 = select i1 %icmp0, i32 %a, i32 %b
@@ -31,20 +87,69 @@ define amdgpu_kernel void @v_sad_u32_pat1(ptr addrspace(1) %out, i32 %a, i32 %b,
}
define amdgpu_kernel void @v_sad_u32_constant_pat1(ptr addrspace(1) %out, i32 %a) {
-; GCN-LABEL: v_sad_u32_constant_pat1:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_load_dword s2, s[8:9], 0x2
-; GCN-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x0
-; GCN-NEXT: s_add_i32 s12, s12, s17
-; GCN-NEXT: v_mov_b32_e32 v0, 0x5a
-; GCN-NEXT: s_mov_b32 flat_scratch_lo, s13
-; GCN-NEXT: s_lshr_b32 flat_scratch_hi, s12, 8
-; GCN-NEXT: s_waitcnt lgkmcnt(0)
-; GCN-NEXT: v_sad_u32 v2, s2, v0, 20
-; GCN-NEXT: v_mov_b32_e32 v0, s0
-; GCN-NEXT: v_mov_b32_e32 v1, s1
-; GCN-NEXT: flat_store_dword v[0:1], v2
-; GCN-NEXT: s_endpgm
+; GFX7-SDAG-LABEL: v_sad_u32_constant_pat1:
+; GFX7-SDAG: ; %bb.0:
+; GFX7-SDAG-NEXT: s_load_dword s6, s[4:5], 0xb
+; GFX7-SDAG-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x9
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v0, 0x5a
+; GFX7-SDAG-NEXT: s_mov_b32 s3, 0xf000
+; GFX7-SDAG-NEXT: s_mov_b32 s2, -1
+; GFX7-SDAG-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-SDAG-NEXT: v_sad_u32 v0, s6, v0, 20
+; GFX7-SDAG-NEXT: buffer_store_dword v0, off, s[0:3], 0
+; GFX7-SDAG-NEXT: s_endpgm
+;
+; GFX7-GISEL-LABEL: v_sad_u32_constant_pat1:
+; GFX7-GISEL: ; %bb.0:
+; GFX7-GISEL-NEXT: s_load_dword s3, s[4:5], 0xb
+; GFX7-GISEL-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x9
+; GFX7-GISEL-NEXT: s_mov_b32 s2, -1
+; GFX7-GISEL-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-GISEL-NEXT: s_max_u32 s4, s3, 0x5a
+; GFX7-GISEL-NEXT: s_min_u32 s3, s3, 0x5a
+; GFX7-GISEL-NEXT: s_sub_i32 s3, s4, s3
+; GFX7-GISEL-NEXT: s_add_i32 s3, s3, 20
+; GFX7-GISEL-NEXT: v_mov_b32_e32 v0, s3
+; GFX7-GISEL-NEXT: s_mov_b32 s3, 0xf000
+; GFX7-GISEL-NEXT: buffer_store_dword v0, off, s[0:3], 0
+; GFX7-GISEL-NEXT: s_endpgm
+;
+; GFX12-5-SDAG-LABEL: v_sad_u32_constant_pat1:
+; GFX12-5-SDAG: ; %bb.0:
+; GFX12-5-SDAG-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-SDAG-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-SDAG-NEXT: v_nop
+; GFX12-5-SDAG-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-SDAG-NEXT: s_load_b96 s[0:2], s[4:5], 0x24 nv
+; GFX12-5-SDAG-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-SDAG-NEXT: s_min_u32 s3, s2, 0x5a
+; GFX12-5-SDAG-NEXT: s_max_u32 s2, s2, 0x5a
+; GFX12-5-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
+; GFX12-5-SDAG-NEXT: s_sub_co_i32 s2, s2, s3
+; GFX12-5-SDAG-NEXT: s_add_co_i32 s2, s2, 20
+; GFX12-5-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
+; GFX12-5-SDAG-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s2
+; GFX12-5-SDAG-NEXT: global_store_b32 v0, v1, s[0:1]
+; GFX12-5-SDAG-NEXT: s_endpgm
+;
+; GFX12-5-GISEL-LABEL: v_sad_u32_constant_pat1:
+; GFX12-5-GISEL: ; %bb.0:
+; GFX12-5-GISEL-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-GISEL-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-GISEL-NEXT: v_nop
+; GFX12-5-GISEL-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-GISEL-NEXT: s_load_b96 s[0:2], s[4:5], 0x24 nv
+; GFX12-5-GISEL-NEXT: v_mov_b32_e32 v1, 0
+; GFX12-5-GISEL-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-GISEL-NEXT: s_max_u32 s3, s2, 0x5a
+; GFX12-5-GISEL-NEXT: s_min_u32 s2, s2, 0x5a
+; GFX12-5-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
+; GFX12-5-GISEL-NEXT: s_sub_co_i32 s2, s3, s2
+; GFX12-5-GISEL-NEXT: s_add_co_i32 s2, s2, 20
+; GFX12-5-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
+; GFX12-5-GISEL-NEXT: v_mov_b32_e32 v0, s2
+; GFX12-5-GISEL-NEXT: global_store_b32 v1, v0, s[0:1]
+; GFX12-5-GISEL-NEXT: s_endpgm
%icmp0 = icmp ugt i32 %a, 90
%t0 = select i1 %icmp0, i32 %a, i32 90
@@ -59,21 +164,77 @@ define amdgpu_kernel void @v_sad_u32_constant_pat1(ptr addrspace(1) %out, i32 %a
}
define amdgpu_kernel void @v_sad_u32_pat2(ptr addrspace(1) %out, i32 %a, i32 %b, i32 %c) {
-; GCN-LABEL: v_sad_u32_pat2:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_load_dwordx4 s[0:3], s[8:9], 0x2
-; GCN-NEXT: s_load_dwordx2 s[4:5], s[8:9], 0x0
-; GCN-NEXT: s_add_i32 s12, s12, s17
-; GCN-NEXT: s_mov_b32 flat_scratch_lo, s13
-; GCN-NEXT: s_lshr_b32 flat_scratch_hi, s12, 8
-; GCN-NEXT: s_waitcnt lgkmcnt(0)
-; GCN-NEXT: v_mov_b32_e32 v0, s0
-; GCN-NEXT: v_mov_b32_e32 v1, s2
-; GCN-NEXT: v_sad_u32 v2, s1, v0, v1
-; GCN-NEXT: v_mov_b32_e32 v0, s4
-; GCN-NEXT: v_mov_b32_e32 v1, s5
-; GCN-NEXT: flat_store_dword v[0:1], v2
-; GCN-NEXT: s_endpgm
+; GFX7-SDAG-LABEL: v_sad_u32_pat2:
+; GFX7-SDAG: ; %bb.0:
+; GFX7-SDAG-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0xb
+; GFX7-SDAG-NEXT: s_load_dwordx2 s[4:5], s[4:5], 0x9
+; GFX7-SDAG-NEXT: s_mov_b32 s7, 0xf000
+; GFX7-SDAG-NEXT: s_mov_b32 s6, -1
+; GFX7-SDAG-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v0, s0
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v1, s2
+; GFX7-SDAG-NEXT: v_sad_u32 v0, s1, v0, v1
+; GFX7-SDAG-NEXT: buffer_store_dword v0, off, s[4:7], 0
+; GFX7-SDAG-NEXT: s_endpgm
+;
+; GFX7-GISEL-LABEL: v_sad_u32_pat2:
+; GFX7-GISEL: ; %bb.0:
+; GFX7-GISEL-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x9
+; GFX7-GISEL-NEXT: s_load_dword s4, s[4:5], 0xd
+; GFX7-GISEL-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-GISEL-NEXT: s_sub_u32 s5, s3, s2
+; GFX7-GISEL-NEXT: s_cselect_b32 s6, 1, 0
+; GFX7-GISEL-NEXT: s_sub_i32 s3, s2, s3
+; GFX7-GISEL-NEXT: s_cmp_lg_u32 s6, 0
+; GFX7-GISEL-NEXT: s_cselect_b32 s3, s3, s5
+; GFX7-GISEL-NEXT: s_add_i32 s3, s3, s4
+; GFX7-GISEL-NEXT: s_mov_b32 s2, -1
+; GFX7-GISEL-NEXT: v_mov_b32_e32 v0, s3
+; GFX7-GISEL-NEXT: s_mov_b32 s3, 0xf000
+; GFX7-GISEL-NEXT: buffer_store_dword v0, off, s[0:3], 0
+; GFX7-GISEL-NEXT: s_endpgm
+;
+; GFX12-5-SDAG-LABEL: v_sad_u32_pat2:
+; GFX12-5-SDAG: ; %bb.0:
+; GFX12-5-SDAG-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-SDAG-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-SDAG-NEXT: v_nop
+; GFX12-5-SDAG-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-SDAG-NEXT: s_clause 0x1
+; GFX12-5-SDAG-NEXT: s_load_b96 s[0:2], s[4:5], 0x2c nv
+; GFX12-5-SDAG-NEXT: s_load_b64 s[6:7], s[4:5], 0x24 nv
+; GFX12-5-SDAG-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-SDAG-NEXT: s_min_u32 s3, s1, s0
+; GFX12-5-SDAG-NEXT: s_max_u32 s0, s1, s0
+; GFX12-5-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
+; GFX12-5-SDAG-NEXT: s_sub_co_i32 s0, s0, s3
+; GFX12-5-SDAG-NEXT: s_add_co_i32 s0, s0, s2
+; GFX12-5-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
+; GFX12-5-SDAG-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s0
+; GFX12-5-SDAG-NEXT: global_store_b32 v0, v1, s[6:7]
+; GFX12-5-SDAG-NEXT: s_endpgm
+;
+; GFX12-5-GISEL-LABEL: v_sad_u32_pat2:
+; GFX12-5-GISEL: ; %bb.0:
+; GFX12-5-GISEL-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-GISEL-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-GISEL-NEXT: v_nop
+; GFX12-5-GISEL-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-GISEL-NEXT: s_clause 0x1
+; GFX12-5-GISEL-NEXT: s_load_b96 s[0:2], s[4:5], 0x2c nv
+; GFX12-5-GISEL-NEXT: s_load_b64 s[6:7], s[4:5], 0x24 nv
+; GFX12-5-GISEL-NEXT: v_mov_b32_e32 v1, 0
+; GFX12-5-GISEL-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-GISEL-NEXT: s_sub_co_u32 s3, s1, s0
+; GFX12-5-GISEL-NEXT: s_cselect_b32 s4, 1, 0
+; GFX12-5-GISEL-NEXT: s_sub_co_i32 s0, s0, s1
+; GFX12-5-GISEL-NEXT: s_cmp_lg_u32 s4, 0
+; GFX12-5-GISEL-NEXT: s_cselect_b32 s0, s0, s3
+; GFX12-5-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
+; GFX12-5-GISEL-NEXT: s_add_co_i32 s0, s0, s2
+; GFX12-5-GISEL-NEXT: v_mov_b32_e32 v0, s0
+; GFX12-5-GISEL-NEXT: global_store_b32 v1, v0, s[6:7]
+; GFX12-5-GISEL-NEXT: s_endpgm
%icmp0 = icmp ugt i32 %a, %b
%sub0 = sub i32 %a, %b
%sub1 = sub i32 %b, %a
@@ -86,30 +247,99 @@ define amdgpu_kernel void @v_sad_u32_pat2(ptr addrspace(1) %out, i32 %a, i32 %b,
}
define amdgpu_kernel void @v_sad_u32_multi_use_sub_pat1(ptr addrspace(1) %out, i32 %a, i32 %b, i32 %c) {
-; GCN-LABEL: v_sad_u32_multi_use_sub_pat1:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_mov_b64 s[22:23], s[2:3]
-; GCN-NEXT: s_mov_b64 s[20:21], s[0:1]
-; GCN-NEXT: s_load_dwordx4 s[0:3], s[8:9], 0x2
-; GCN-NEXT: s_load_dwordx2 s[4:5], s[8:9], 0x0
-; GCN-NEXT: s_add_i32 s12, s12, s17
-; GCN-NEXT: s_lshr_b32 flat_scratch_hi, s12, 8
-; GCN-NEXT: s_add_u32 s20, s20, s17
-; GCN-NEXT: s_addc_u32 s21, s21, 0
-; GCN-NEXT: s_waitcnt lgkmcnt(0)
-; GCN-NEXT: s_min_u32 s3, s0, s1
-; GCN-NEXT: s_max_u32 s0, s0, s1
-; GCN-NEXT: s_sub_i32 s0, s0, s3
-; GCN-NEXT: v_mov_b32_e32 v0, s4
-; GCN-NEXT: v_mov_b32_e32 v2, s0
-; GCN-NEXT: s_add_i32 s0, s0, s2
-; GCN-NEXT: s_mov_b32 flat_scratch_lo, s13
-; GCN-NEXT: v_mov_b32_e32 v1, s5
-; GCN-NEXT: buffer_store_dword v2, v0, s[20:23], 0 offen
-; GCN-NEXT: s_waitcnt vmcnt(0)
-; GCN-NEXT: v_mov_b32_e32 v2, s0
-; GCN-NEXT: flat_store_dword v[0:1], v2
-; GCN-NEXT: s_endpgm
+; GFX7-SDAG-LABEL: v_sad_u32_multi_use_sub_pat1:
+; GFX7-SDAG: ; %bb.0:
+; GFX7-SDAG-NEXT: s_mov_b32 s12, SCRATCH_RSRC_DWORD0
+; GFX7-SDAG-NEXT: s_mov_b32 s13, SCRATCH_RSRC_DWORD1
+; GFX7-SDAG-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0xb
+; GFX7-SDAG-NEXT: s_load_dwordx2 s[4:5], s[4:5], 0x9
+; GFX7-SDAG-NEXT: s_mov_b32 s14, -1
+; GFX7-SDAG-NEXT: s_mov_b32 s15, 0xe8f000
+; GFX7-SDAG-NEXT: s_add_u32 s12, s12, s11
+; GFX7-SDAG-NEXT: s_addc_u32 s13, s13, 0
+; GFX7-SDAG-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-SDAG-NEXT: s_min_u32 s3, s0, s1
+; GFX7-SDAG-NEXT: s_max_u32 s0, s0, s1
+; GFX7-SDAG-NEXT: s_sub_i32 s0, s0, s3
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v0, s0
+; GFX7-SDAG-NEXT: s_add_i32 s0, s0, s2
+; GFX7-SDAG-NEXT: s_mov_b32 s7, 0xf000
+; GFX7-SDAG-NEXT: s_mov_b32 s6, -1
+; GFX7-SDAG-NEXT: buffer_store_dword v0, v0, s[12:15], 0 offen
+; GFX7-SDAG-NEXT: s_waitcnt vmcnt(0)
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v0, s0
+; GFX7-SDAG-NEXT: buffer_store_dword v0, off, s[4:7], 0
+; GFX7-SDAG-NEXT: s_endpgm
+;
+; GFX7-GISEL-LABEL: v_sad_u32_multi_use_sub_pat1:
+; GFX7-GISEL: ; %bb.0:
+; GFX7-GISEL-NEXT: s_mov_b32 s12, SCRATCH_RSRC_DWORD0
+; GFX7-GISEL-NEXT: s_mov_b32 s13, SCRATCH_RSRC_DWORD1
+; GFX7-GISEL-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x9
+; GFX7-GISEL-NEXT: s_load_dword s4, s[4:5], 0xd
+; GFX7-GISEL-NEXT: s_mov_b32 s14, -1
+; GFX7-GISEL-NEXT: s_mov_b32 s15, 0xe8f000
+; GFX7-GISEL-NEXT: s_add_u32 s12, s12, s11
+; GFX7-GISEL-NEXT: s_addc_u32 s13, s13, 0
+; GFX7-GISEL-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-GISEL-NEXT: s_max_u32 s5, s2, s3
+; GFX7-GISEL-NEXT: s_min_u32 s2, s2, s3
+; GFX7-GISEL-NEXT: s_sub_i32 s2, s5, s2
+; GFX7-GISEL-NEXT: v_mov_b32_e32 v0, s2
+; GFX7-GISEL-NEXT: s_add_i32 s2, s2, s4
+; GFX7-GISEL-NEXT: buffer_store_dword v0, v0, s[12:15], 0 offen
+; GFX7-GISEL-NEXT: s_waitcnt vmcnt(0)
+; GFX7-GISEL-NEXT: v_mov_b32_e32 v0, s2
+; GFX7-GISEL-NEXT: s_mov_b32 s2, -1
+; GFX7-GISEL-NEXT: s_mov_b32 s3, 0xf000
+; GFX7-GISEL-NEXT: buffer_store_dword v0, off, s[0:3], 0
+; GFX7-GISEL-NEXT: s_endpgm
+;
+; GFX12-5-SDAG-LABEL: v_sad_u32_multi_use_sub_pat1:
+; GFX12-5-SDAG: ; %bb.0:
+; GFX12-5-SDAG-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-SDAG-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-SDAG-NEXT: v_nop
+; GFX12-5-SDAG-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-SDAG-NEXT: s_clause 0x1
+; GFX12-5-SDAG-NEXT: s_load_b96 s[0:2], s[4:5], 0x2c nv
+; GFX12-5-SDAG-NEXT: s_load_b64 s[6:7], s[4:5], 0x24 nv
+; GFX12-5-SDAG-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-SDAG-NEXT: s_min_u32 s3, s0, s1
+; GFX12-5-SDAG-NEXT: s_max_u32 s0, s0, s1
+; GFX12-5-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
+; GFX12-5-SDAG-NEXT: s_sub_co_i32 s0, s0, s3
+; GFX12-5-SDAG-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s0
+; GFX12-5-SDAG-NEXT: s_add_co_i32 s1, s0, s2
+; GFX12-5-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
+; GFX12-5-SDAG-NEXT: v_mov_b32_e32 v2, s1
+; GFX12-5-SDAG-NEXT: scratch_store_b32 off, v1, s0 scope:SCOPE_SYS
+; GFX12-5-SDAG-NEXT: s_wait_storecnt 0x0
+; GFX12-5-SDAG-NEXT: global_store_b32 v0, v2, s[6:7]
+; GFX12-5-SDAG-NEXT: s_endpgm
+;
+; GFX12-5-GISEL-LABEL: v_sad_u32_multi_use_sub_pat1:
+; GFX12-5-GISEL: ; %bb.0:
+; GFX12-5-GISEL-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-GISEL-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-GISEL-NEXT: v_nop
+; GFX12-5-GISEL-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-GISEL-NEXT: s_clause 0x1
+; GFX12-5-GISEL-NEXT: s_load_b96 s[0:2], s[4:5], 0x2c nv
+; GFX12-5-GISEL-NEXT: s_load_b64 s[6:7], s[4:5], 0x24 nv
+; GFX12-5-GISEL-NEXT: v_mov_b32_e32 v2, 0
+; GFX12-5-GISEL-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-GISEL-NEXT: s_max_u32 s3, s0, s1
+; GFX12-5-GISEL-NEXT: s_min_u32 s0, s0, s1
+; GFX12-5-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
+; GFX12-5-GISEL-NEXT: s_sub_co_i32 s0, s3, s0
+; GFX12-5-GISEL-NEXT: s_add_co_i32 s1, s0, s2
+; GFX12-5-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
+; GFX12-5-GISEL-NEXT: v_dual_mov_b32 v0, s0 :: v_dual_mov_b32 v1, s1
+; GFX12-5-GISEL-NEXT: scratch_store_b32 off, v0, s0 scope:SCOPE_SYS
+; GFX12-5-GISEL-NEXT: s_wait_storecnt 0x0
+; GFX12-5-GISEL-NEXT: global_store_b32 v2, v1, s[6:7]
+; GFX12-5-GISEL-NEXT: s_endpgm
%icmp0 = icmp ugt i32 %a, %b
%t0 = select i1 %icmp0, i32 %a, i32 %b
@@ -125,27 +355,95 @@ define amdgpu_kernel void @v_sad_u32_multi_use_sub_pat1(ptr addrspace(1) %out, i
}
define amdgpu_kernel void @v_sad_u32_multi_use_add_pat1(ptr addrspace(1) %out, i32 %a, i32 %b, i32 %c) {
-; GCN-LABEL: v_sad_u32_multi_use_add_pat1:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_mov_b64 s[22:23], s[2:3]
-; GCN-NEXT: s_mov_b64 s[20:21], s[0:1]
-; GCN-NEXT: s_load_dwordx4 s[0:3], s[8:9], 0x2
-; GCN-NEXT: s_load_dwordx2 s[4:5], s[8:9], 0x0
-; GCN-NEXT: s_add_i32 s12, s12, s17
-; GCN-NEXT: s_lshr_b32 flat_scratch_hi, s12, 8
-; GCN-NEXT: s_add_u32 s20, s20, s17
-; GCN-NEXT: s_waitcnt lgkmcnt(0)
-; GCN-NEXT: v_mov_b32_e32 v2, s1
-; GCN-NEXT: v_mov_b32_e32 v3, s2
-; GCN-NEXT: s_mov_b32 flat_scratch_lo, s13
-; GCN-NEXT: s_addc_u32 s21, s21, 0
-; GCN-NEXT: v_mov_b32_e32 v0, s4
-; GCN-NEXT: v_mov_b32_e32 v1, s5
-; GCN-NEXT: v_sad_u32 v2, s0, v2, v3
-; GCN-NEXT: buffer_store_dword v2, v0, s[20:23], 0 offen
-; GCN-NEXT: s_waitcnt vmcnt(0)
-; GCN-NEXT: flat_store_dword v[0:1], v2
-; GCN-NEXT: s_endpgm
+; GFX7-SDAG-LABEL: v_sad_u32_multi_use_add_pat1:
+; GFX7-SDAG: ; %bb.0:
+; GFX7-SDAG-NEXT: s_mov_b32 s12, SCRATCH_RSRC_DWORD0
+; GFX7-SDAG-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0xb
+; GFX7-SDAG-NEXT: s_load_dwordx2 s[4:5], s[4:5], 0x9
+; GFX7-SDAG-NEXT: s_mov_b32 s13, SCRATCH_RSRC_DWORD1
+; GFX7-SDAG-NEXT: s_mov_b32 s14, -1
+; GFX7-SDAG-NEXT: s_mov_b32 s15, 0xe8f000
+; GFX7-SDAG-NEXT: s_add_u32 s12, s12, s11
+; GFX7-SDAG-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v0, s1
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v1, s2
+; GFX7-SDAG-NEXT: s_addc_u32 s13, s13, 0
+; GFX7-SDAG-NEXT: v_sad_u32 v0, s0, v0, v1
+; GFX7-SDAG-NEXT: s_mov_b32 s7, 0xf000
+; GFX7-SDAG-NEXT: s_mov_b32 s6, -1
+; GFX7-SDAG-NEXT: buffer_store_dword v0, v0, s[12:15], 0 offen
+; GFX7-SDAG-NEXT: s_waitcnt vmcnt(0)
+; GFX7-SDAG-NEXT: buffer_store_dword v0, off, s[4:7], 0
+; GFX7-SDAG-NEXT: s_endpgm
+;
+; GFX7-GISEL-LABEL: v_sad_u32_multi_use_add_pat1:
+; GFX7-GISEL: ; %bb.0:
+; GFX7-GISEL-NEXT: s_mov_b32 s12, SCRATCH_RSRC_DWORD0
+; GFX7-GISEL-NEXT: s_mov_b32 s13, SCRATCH_RSRC_DWORD1
+; GFX7-GISEL-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x9
+; GFX7-GISEL-NEXT: s_load_dword s4, s[4:5], 0xd
+; GFX7-GISEL-NEXT: s_mov_b32 s14, -1
+; GFX7-GISEL-NEXT: s_mov_b32 s15, 0xe8f000
+; GFX7-GISEL-NEXT: s_add_u32 s12, s12, s11
+; GFX7-GISEL-NEXT: s_addc_u32 s13, s13, 0
+; GFX7-GISEL-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-GISEL-NEXT: s_max_u32 s5, s2, s3
+; GFX7-GISEL-NEXT: s_min_u32 s6, s2, s3
+; GFX7-GISEL-NEXT: s_sub_i32 s5, s5, s6
+; GFX7-GISEL-NEXT: s_add_i32 s4, s5, s4
+; GFX7-GISEL-NEXT: v_mov_b32_e32 v0, s4
+; GFX7-GISEL-NEXT: s_mov_b32 s2, -1
+; GFX7-GISEL-NEXT: s_mov_b32 s3, 0xf000
+; GFX7-GISEL-NEXT: buffer_store_dword v0, v0, s[12:15], 0 offen
+; GFX7-GISEL-NEXT: s_waitcnt vmcnt(0)
+; GFX7-GISEL-NEXT: buffer_store_dword v0, off, s[0:3], 0
+; GFX7-GISEL-NEXT: s_endpgm
+;
+; GFX12-5-SDAG-LABEL: v_sad_u32_multi_use_add_pat1:
+; GFX12-5-SDAG: ; %bb.0:
+; GFX12-5-SDAG-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-SDAG-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-SDAG-NEXT: v_nop
+; GFX12-5-SDAG-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-SDAG-NEXT: s_clause 0x1
+; GFX12-5-SDAG-NEXT: s_load_b96 s[0:2], s[4:5], 0x2c nv
+; GFX12-5-SDAG-NEXT: s_load_b64 s[6:7], s[4:5], 0x24 nv
+; GFX12-5-SDAG-NEXT: v_mov_b32_e32 v1, 0
+; GFX12-5-SDAG-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-SDAG-NEXT: s_min_u32 s3, s0, s1
+; GFX12-5-SDAG-NEXT: s_max_u32 s0, s0, s1
+; GFX12-5-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
+; GFX12-5-SDAG-NEXT: s_sub_co_i32 s0, s0, s3
+; GFX12-5-SDAG-NEXT: s_add_co_i32 s0, s0, s2
+; GFX12-5-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
+; GFX12-5-SDAG-NEXT: v_mov_b32_e32 v0, s0
+; GFX12-5-SDAG-NEXT: scratch_store_b32 off, v0, s0 scope:SCOPE_SYS
+; GFX12-5-SDAG-NEXT: s_wait_storecnt 0x0
+; GFX12-5-SDAG-NEXT: global_store_b32 v1, v0, s[6:7]
+; GFX12-5-SDAG-NEXT: s_endpgm
+;
+; GFX12-5-GISEL-LABEL: v_sad_u32_multi_use_add_pat1:
+; GFX12-5-GISEL: ; %bb.0:
+; GFX12-5-GISEL-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-GISEL-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-GISEL-NEXT: v_nop
+; GFX12-5-GISEL-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-GISEL-NEXT: s_clause 0x1
+; GFX12-5-GISEL-NEXT: s_load_b96 s[0:2], s[4:5], 0x2c nv
+; GFX12-5-GISEL-NEXT: s_load_b64 s[6:7], s[4:5], 0x24 nv
+; GFX12-5-GISEL-NEXT: v_mov_b32_e32 v1, 0
+; GFX12-5-GISEL-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-GISEL-NEXT: s_max_u32 s3, s0, s1
+; GFX12-5-GISEL-NEXT: s_min_u32 s0, s0, s1
+; GFX12-5-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
+; GFX12-5-GISEL-NEXT: s_sub_co_i32 s0, s3, s0
+; GFX12-5-GISEL-NEXT: s_add_co_i32 s0, s0, s2
+; GFX12-5-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
+; GFX12-5-GISEL-NEXT: v_mov_b32_e32 v0, s0
+; GFX12-5-GISEL-NEXT: scratch_store_b32 off, v0, s0 scope:SCOPE_SYS
+; GFX12-5-GISEL-NEXT: s_wait_storecnt 0x0
+; GFX12-5-GISEL-NEXT: global_store_b32 v1, v0, s[6:7]
+; GFX12-5-GISEL-NEXT: s_endpgm
%icmp0 = icmp ugt i32 %a, %b
%t0 = select i1 %icmp0, i32 %a, i32 %b
@@ -160,29 +458,97 @@ define amdgpu_kernel void @v_sad_u32_multi_use_add_pat1(ptr addrspace(1) %out, i
}
define amdgpu_kernel void @v_sad_u32_multi_use_max_pat1(ptr addrspace(1) %out, i32 %a, i32 %b, i32 %c) {
-; GCN-LABEL: v_sad_u32_multi_use_max_pat1:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_mov_b64 s[22:23], s[2:3]
-; GCN-NEXT: s_mov_b64 s[20:21], s[0:1]
-; GCN-NEXT: s_load_dwordx4 s[0:3], s[8:9], 0x2
-; GCN-NEXT: s_load_dwordx2 s[4:5], s[8:9], 0x0
-; GCN-NEXT: s_add_i32 s12, s12, s17
-; GCN-NEXT: s_lshr_b32 flat_scratch_hi, s12, 8
-; GCN-NEXT: s_add_u32 s20, s20, s17
-; GCN-NEXT: s_addc_u32 s21, s21, 0
-; GCN-NEXT: s_waitcnt lgkmcnt(0)
-; GCN-NEXT: s_max_u32 s3, s0, s1
-; GCN-NEXT: v_mov_b32_e32 v0, s1
-; GCN-NEXT: v_mov_b32_e32 v1, s2
-; GCN-NEXT: v_mov_b32_e32 v2, s3
-; GCN-NEXT: s_mov_b32 flat_scratch_lo, s13
-; GCN-NEXT: v_sad_u32 v3, s0, v0, v1
-; GCN-NEXT: buffer_store_dword v2, v0, s[20:23], 0 offen
-; GCN-NEXT: s_waitcnt vmcnt(0)
-; GCN-NEXT: v_mov_b32_e32 v0, s4
-; GCN-NEXT: v_mov_b32_e32 v1, s5
-; GCN-NEXT: flat_store_dword v[0:1], v3
-; GCN-NEXT: s_endpgm
+; GFX7-SDAG-LABEL: v_sad_u32_multi_use_max_pat1:
+; GFX7-SDAG: ; %bb.0:
+; GFX7-SDAG-NEXT: s_mov_b32 s12, SCRATCH_RSRC_DWORD0
+; GFX7-SDAG-NEXT: s_mov_b32 s13, SCRATCH_RSRC_DWORD1
+; GFX7-SDAG-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0xb
+; GFX7-SDAG-NEXT: s_load_dwordx2 s[4:5], s[4:5], 0x9
+; GFX7-SDAG-NEXT: s_mov_b32 s14, -1
+; GFX7-SDAG-NEXT: s_mov_b32 s15, 0xe8f000
+; GFX7-SDAG-NEXT: s_add_u32 s12, s12, s11
+; GFX7-SDAG-NEXT: s_addc_u32 s13, s13, 0
+; GFX7-SDAG-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-SDAG-NEXT: s_max_u32 s3, s0, s1
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v0, s1
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v1, s3
+; GFX7-SDAG-NEXT: buffer_store_dword v1, v0, s[12:15], 0 offen
+; GFX7-SDAG-NEXT: s_waitcnt vmcnt(0)
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v1, s2
+; GFX7-SDAG-NEXT: s_mov_b32 s7, 0xf000
+; GFX7-SDAG-NEXT: s_mov_b32 s6, -1
+; GFX7-SDAG-NEXT: v_sad_u32 v0, s0, v0, v1
+; GFX7-SDAG-NEXT: buffer_store_dword v0, off, s[4:7], 0
+; GFX7-SDAG-NEXT: s_endpgm
+;
+; GFX7-GISEL-LABEL: v_sad_u32_multi_use_max_pat1:
+; GFX7-GISEL: ; %bb.0:
+; GFX7-GISEL-NEXT: s_mov_b32 s12, SCRATCH_RSRC_DWORD0
+; GFX7-GISEL-NEXT: s_mov_b32 s13, SCRATCH_RSRC_DWORD1
+; GFX7-GISEL-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x9
+; GFX7-GISEL-NEXT: s_load_dword s4, s[4:5], 0xd
+; GFX7-GISEL-NEXT: s_mov_b32 s14, -1
+; GFX7-GISEL-NEXT: s_mov_b32 s15, 0xe8f000
+; GFX7-GISEL-NEXT: s_add_u32 s12, s12, s11
+; GFX7-GISEL-NEXT: s_addc_u32 s13, s13, 0
+; GFX7-GISEL-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-GISEL-NEXT: s_max_u32 s5, s2, s3
+; GFX7-GISEL-NEXT: s_min_u32 s2, s2, s3
+; GFX7-GISEL-NEXT: s_sub_i32 s2, s5, s2
+; GFX7-GISEL-NEXT: v_mov_b32_e32 v0, s5
+; GFX7-GISEL-NEXT: s_add_i32 s2, s2, s4
+; GFX7-GISEL-NEXT: buffer_store_dword v0, v0, s[12:15], 0 offen
+; GFX7-GISEL-NEXT: s_waitcnt vmcnt(0)
+; GFX7-GISEL-NEXT: v_mov_b32_e32 v0, s2
+; GFX7-GISEL-NEXT: s_mov_b32 s2, -1
+; GFX7-GISEL-NEXT: s_mov_b32 s3, 0xf000
+; GFX7-GISEL-NEXT: buffer_store_dword v0, off, s[0:3], 0
+; GFX7-GISEL-NEXT: s_endpgm
+;
+; GFX12-5-SDAG-LABEL: v_sad_u32_multi_use_max_pat1:
+; GFX12-5-SDAG: ; %bb.0:
+; GFX12-5-SDAG-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-SDAG-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-SDAG-NEXT: v_nop
+; GFX12-5-SDAG-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-SDAG-NEXT: s_clause 0x1
+; GFX12-5-SDAG-NEXT: s_load_b96 s[0:2], s[4:5], 0x2c nv
+; GFX12-5-SDAG-NEXT: s_load_b64 s[6:7], s[4:5], 0x24 nv
+; GFX12-5-SDAG-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-SDAG-NEXT: s_max_u32 s3, s0, s1
+; GFX12-5-SDAG-NEXT: s_min_u32 s0, s0, s1
+; GFX12-5-SDAG-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s3
+; GFX12-5-SDAG-NEXT: s_sub_co_i32 s0, s3, s0
+; GFX12-5-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
+; GFX12-5-SDAG-NEXT: s_add_co_i32 s0, s0, s2
+; GFX12-5-SDAG-NEXT: v_mov_b32_e32 v2, s0
+; GFX12-5-SDAG-NEXT: scratch_store_b32 off, v1, s0 scope:SCOPE_SYS
+; GFX12-5-SDAG-NEXT: s_wait_storecnt 0x0
+; GFX12-5-SDAG-NEXT: global_store_b32 v0, v2, s[6:7]
+; GFX12-5-SDAG-NEXT: s_endpgm
+;
+; GFX12-5-GISEL-LABEL: v_sad_u32_multi_use_max_pat1:
+; GFX12-5-GISEL: ; %bb.0:
+; GFX12-5-GISEL-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-GISEL-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-GISEL-NEXT: v_nop
+; GFX12-5-GISEL-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-GISEL-NEXT: s_clause 0x1
+; GFX12-5-GISEL-NEXT: s_load_b96 s[0:2], s[4:5], 0x2c nv
+; GFX12-5-GISEL-NEXT: s_load_b64 s[6:7], s[4:5], 0x24 nv
+; GFX12-5-GISEL-NEXT: v_mov_b32_e32 v2, 0
+; GFX12-5-GISEL-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-GISEL-NEXT: s_max_u32 s3, s0, s1
+; GFX12-5-GISEL-NEXT: s_min_u32 s0, s0, s1
+; GFX12-5-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
+; GFX12-5-GISEL-NEXT: s_sub_co_i32 s0, s3, s0
+; GFX12-5-GISEL-NEXT: s_add_co_i32 s0, s0, s2
+; GFX12-5-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
+; GFX12-5-GISEL-NEXT: v_dual_mov_b32 v0, s3 :: v_dual_mov_b32 v1, s0
+; GFX12-5-GISEL-NEXT: scratch_store_b32 off, v0, s0 scope:SCOPE_SYS
+; GFX12-5-GISEL-NEXT: s_wait_storecnt 0x0
+; GFX12-5-GISEL-NEXT: global_store_b32 v2, v1, s[6:7]
+; GFX12-5-GISEL-NEXT: s_endpgm
%icmp0 = icmp ugt i32 %a, %b
%t0 = select i1 %icmp0, i32 %a, i32 %b
store volatile i32 %t0, ptr addrspace(5) poison
@@ -198,29 +564,97 @@ define amdgpu_kernel void @v_sad_u32_multi_use_max_pat1(ptr addrspace(1) %out, i
}
define amdgpu_kernel void @v_sad_u32_multi_use_min_pat1(ptr addrspace(1) %out, i32 %a, i32 %b, i32 %c) {
-; GCN-LABEL: v_sad_u32_multi_use_min_pat1:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_mov_b64 s[22:23], s[2:3]
-; GCN-NEXT: s_mov_b64 s[20:21], s[0:1]
-; GCN-NEXT: s_load_dwordx4 s[0:3], s[8:9], 0x2
-; GCN-NEXT: s_load_dwordx2 s[4:5], s[8:9], 0x0
-; GCN-NEXT: s_add_i32 s12, s12, s17
-; GCN-NEXT: s_lshr_b32 flat_scratch_hi, s12, 8
-; GCN-NEXT: s_add_u32 s20, s20, s17
-; GCN-NEXT: s_addc_u32 s21, s21, 0
-; GCN-NEXT: s_waitcnt lgkmcnt(0)
-; GCN-NEXT: s_min_u32 s3, s0, s1
-; GCN-NEXT: v_mov_b32_e32 v0, s1
-; GCN-NEXT: v_mov_b32_e32 v1, s2
-; GCN-NEXT: v_mov_b32_e32 v2, s3
-; GCN-NEXT: s_mov_b32 flat_scratch_lo, s13
-; GCN-NEXT: v_sad_u32 v3, s0, v0, v1
-; GCN-NEXT: buffer_store_dword v2, v0, s[20:23], 0 offen
-; GCN-NEXT: s_waitcnt vmcnt(0)
-; GCN-NEXT: v_mov_b32_e32 v0, s4
-; GCN-NEXT: v_mov_b32_e32 v1, s5
-; GCN-NEXT: flat_store_dword v[0:1], v3
-; GCN-NEXT: s_endpgm
+; GFX7-SDAG-LABEL: v_sad_u32_multi_use_min_pat1:
+; GFX7-SDAG: ; %bb.0:
+; GFX7-SDAG-NEXT: s_mov_b32 s12, SCRATCH_RSRC_DWORD0
+; GFX7-SDAG-NEXT: s_mov_b32 s13, SCRATCH_RSRC_DWORD1
+; GFX7-SDAG-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0xb
+; GFX7-SDAG-NEXT: s_load_dwordx2 s[4:5], s[4:5], 0x9
+; GFX7-SDAG-NEXT: s_mov_b32 s14, -1
+; GFX7-SDAG-NEXT: s_mov_b32 s15, 0xe8f000
+; GFX7-SDAG-NEXT: s_add_u32 s12, s12, s11
+; GFX7-SDAG-NEXT: s_addc_u32 s13, s13, 0
+; GFX7-SDAG-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-SDAG-NEXT: s_min_u32 s3, s0, s1
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v0, s1
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v1, s3
+; GFX7-SDAG-NEXT: buffer_store_dword v1, v0, s[12:15], 0 offen
+; GFX7-SDAG-NEXT: s_waitcnt vmcnt(0)
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v1, s2
+; GFX7-SDAG-NEXT: s_mov_b32 s7, 0xf000
+; GFX7-SDAG-NEXT: s_mov_b32 s6, -1
+; GFX7-SDAG-NEXT: v_sad_u32 v0, s0, v0, v1
+; GFX7-SDAG-NEXT: buffer_store_dword v0, off, s[4:7], 0
+; GFX7-SDAG-NEXT: s_endpgm
+;
+; GFX7-GISEL-LABEL: v_sad_u32_multi_use_min_pat1:
+; GFX7-GISEL: ; %bb.0:
+; GFX7-GISEL-NEXT: s_mov_b32 s12, SCRATCH_RSRC_DWORD0
+; GFX7-GISEL-NEXT: s_mov_b32 s13, SCRATCH_RSRC_DWORD1
+; GFX7-GISEL-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x9
+; GFX7-GISEL-NEXT: s_load_dword s4, s[4:5], 0xd
+; GFX7-GISEL-NEXT: s_mov_b32 s14, -1
+; GFX7-GISEL-NEXT: s_mov_b32 s15, 0xe8f000
+; GFX7-GISEL-NEXT: s_add_u32 s12, s12, s11
+; GFX7-GISEL-NEXT: s_addc_u32 s13, s13, 0
+; GFX7-GISEL-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-GISEL-NEXT: s_max_u32 s5, s2, s3
+; GFX7-GISEL-NEXT: s_min_u32 s2, s2, s3
+; GFX7-GISEL-NEXT: v_mov_b32_e32 v0, s2
+; GFX7-GISEL-NEXT: s_sub_i32 s2, s5, s2
+; GFX7-GISEL-NEXT: s_add_i32 s2, s2, s4
+; GFX7-GISEL-NEXT: buffer_store_dword v0, v0, s[12:15], 0 offen
+; GFX7-GISEL-NEXT: s_waitcnt vmcnt(0)
+; GFX7-GISEL-NEXT: v_mov_b32_e32 v0, s2
+; GFX7-GISEL-NEXT: s_mov_b32 s2, -1
+; GFX7-GISEL-NEXT: s_mov_b32 s3, 0xf000
+; GFX7-GISEL-NEXT: buffer_store_dword v0, off, s[0:3], 0
+; GFX7-GISEL-NEXT: s_endpgm
+;
+; GFX12-5-SDAG-LABEL: v_sad_u32_multi_use_min_pat1:
+; GFX12-5-SDAG: ; %bb.0:
+; GFX12-5-SDAG-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-SDAG-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-SDAG-NEXT: v_nop
+; GFX12-5-SDAG-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-SDAG-NEXT: s_clause 0x1
+; GFX12-5-SDAG-NEXT: s_load_b96 s[0:2], s[4:5], 0x2c nv
+; GFX12-5-SDAG-NEXT: s_load_b64 s[6:7], s[4:5], 0x24 nv
+; GFX12-5-SDAG-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-SDAG-NEXT: s_min_u32 s3, s0, s1
+; GFX12-5-SDAG-NEXT: s_max_u32 s0, s0, s1
+; GFX12-5-SDAG-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s3
+; GFX12-5-SDAG-NEXT: s_sub_co_i32 s0, s0, s3
+; GFX12-5-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
+; GFX12-5-SDAG-NEXT: s_add_co_i32 s0, s0, s2
+; GFX12-5-SDAG-NEXT: v_mov_b32_e32 v2, s0
+; GFX12-5-SDAG-NEXT: scratch_store_b32 off, v1, s0 scope:SCOPE_SYS
+; GFX12-5-SDAG-NEXT: s_wait_storecnt 0x0
+; GFX12-5-SDAG-NEXT: global_store_b32 v0, v2, s[6:7]
+; GFX12-5-SDAG-NEXT: s_endpgm
+;
+; GFX12-5-GISEL-LABEL: v_sad_u32_multi_use_min_pat1:
+; GFX12-5-GISEL: ; %bb.0:
+; GFX12-5-GISEL-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-GISEL-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-GISEL-NEXT: v_nop
+; GFX12-5-GISEL-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-GISEL-NEXT: s_clause 0x1
+; GFX12-5-GISEL-NEXT: s_load_b96 s[0:2], s[4:5], 0x2c nv
+; GFX12-5-GISEL-NEXT: s_load_b64 s[6:7], s[4:5], 0x24 nv
+; GFX12-5-GISEL-NEXT: v_mov_b32_e32 v2, 0
+; GFX12-5-GISEL-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-GISEL-NEXT: s_max_u32 s3, s0, s1
+; GFX12-5-GISEL-NEXT: s_min_u32 s0, s0, s1
+; GFX12-5-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
+; GFX12-5-GISEL-NEXT: s_sub_co_i32 s1, s3, s0
+; GFX12-5-GISEL-NEXT: s_add_co_i32 s1, s1, s2
+; GFX12-5-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
+; GFX12-5-GISEL-NEXT: v_dual_mov_b32 v0, s0 :: v_dual_mov_b32 v1, s1
+; GFX12-5-GISEL-NEXT: scratch_store_b32 off, v0, s0 scope:SCOPE_SYS
+; GFX12-5-GISEL-NEXT: s_wait_storecnt 0x0
+; GFX12-5-GISEL-NEXT: global_store_b32 v2, v1, s[6:7]
+; GFX12-5-GISEL-NEXT: s_endpgm
%icmp0 = icmp ugt i32 %a, %b
%t0 = select i1 %icmp0, i32 %a, i32 %b
@@ -237,29 +671,101 @@ define amdgpu_kernel void @v_sad_u32_multi_use_min_pat1(ptr addrspace(1) %out, i
}
define amdgpu_kernel void @v_sad_u32_multi_use_sub_pat2(ptr addrspace(1) %out, i32 %a, i32 %b, i32 %c) {
-; GCN-LABEL: v_sad_u32_multi_use_sub_pat2:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_mov_b64 s[22:23], s[2:3]
-; GCN-NEXT: s_mov_b64 s[20:21], s[0:1]
-; GCN-NEXT: s_load_dwordx4 s[0:3], s[8:9], 0x2
-; GCN-NEXT: s_load_dwordx2 s[4:5], s[8:9], 0x0
-; GCN-NEXT: s_add_i32 s12, s12, s17
-; GCN-NEXT: s_lshr_b32 flat_scratch_hi, s12, 8
-; GCN-NEXT: s_add_u32 s20, s20, s17
-; GCN-NEXT: s_addc_u32 s21, s21, 0
-; GCN-NEXT: s_waitcnt lgkmcnt(0)
-; GCN-NEXT: s_sub_i32 s3, s0, s1
-; GCN-NEXT: v_mov_b32_e32 v0, s0
-; GCN-NEXT: v_mov_b32_e32 v1, s2
-; GCN-NEXT: v_mov_b32_e32 v2, s3
-; GCN-NEXT: s_mov_b32 flat_scratch_lo, s13
-; GCN-NEXT: v_sad_u32 v3, s1, v0, v1
-; GCN-NEXT: buffer_store_dword v2, v0, s[20:23], 0 offen
-; GCN-NEXT: s_waitcnt vmcnt(0)
-; GCN-NEXT: v_mov_b32_e32 v0, s4
-; GCN-NEXT: v_mov_b32_e32 v1, s5
-; GCN-NEXT: flat_store_dword v[0:1], v3
-; GCN-NEXT: s_endpgm
+; GFX7-SDAG-LABEL: v_sad_u32_multi_use_sub_pat2:
+; GFX7-SDAG: ; %bb.0:
+; GFX7-SDAG-NEXT: s_mov_b32 s12, SCRATCH_RSRC_DWORD0
+; GFX7-SDAG-NEXT: s_mov_b32 s13, SCRATCH_RSRC_DWORD1
+; GFX7-SDAG-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0xb
+; GFX7-SDAG-NEXT: s_load_dwordx2 s[4:5], s[4:5], 0x9
+; GFX7-SDAG-NEXT: s_mov_b32 s14, -1
+; GFX7-SDAG-NEXT: s_mov_b32 s15, 0xe8f000
+; GFX7-SDAG-NEXT: s_add_u32 s12, s12, s11
+; GFX7-SDAG-NEXT: s_addc_u32 s13, s13, 0
+; GFX7-SDAG-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-SDAG-NEXT: s_sub_i32 s3, s0, s1
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v0, s0
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v1, s3
+; GFX7-SDAG-NEXT: buffer_store_dword v1, v0, s[12:15], 0 offen
+; GFX7-SDAG-NEXT: s_waitcnt vmcnt(0)
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v1, s2
+; GFX7-SDAG-NEXT: s_mov_b32 s7, 0xf000
+; GFX7-SDAG-NEXT: s_mov_b32 s6, -1
+; GFX7-SDAG-NEXT: v_sad_u32 v0, s1, v0, v1
+; GFX7-SDAG-NEXT: buffer_store_dword v0, off, s[4:7], 0
+; GFX7-SDAG-NEXT: s_endpgm
+;
+; GFX7-GISEL-LABEL: v_sad_u32_multi_use_sub_pat2:
+; GFX7-GISEL: ; %bb.0:
+; GFX7-GISEL-NEXT: s_mov_b32 s12, SCRATCH_RSRC_DWORD0
+; GFX7-GISEL-NEXT: s_mov_b32 s13, SCRATCH_RSRC_DWORD1
+; GFX7-GISEL-NEXT: s_mov_b32 s14, -1
+; GFX7-GISEL-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x9
+; GFX7-GISEL-NEXT: s_load_dword s4, s[4:5], 0xd
+; GFX7-GISEL-NEXT: s_mov_b32 s15, 0xe8f000
+; GFX7-GISEL-NEXT: s_add_u32 s12, s12, s11
+; GFX7-GISEL-NEXT: s_addc_u32 s13, s13, 0
+; GFX7-GISEL-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-GISEL-NEXT: s_sub_u32 s5, s3, s2
+; GFX7-GISEL-NEXT: s_cselect_b32 s6, 1, 0
+; GFX7-GISEL-NEXT: s_sub_i32 s2, s2, s3
+; GFX7-GISEL-NEXT: s_cmp_lg_u32 s6, 0
+; GFX7-GISEL-NEXT: v_mov_b32_e32 v0, s2
+; GFX7-GISEL-NEXT: s_cselect_b32 s2, s2, s5
+; GFX7-GISEL-NEXT: s_add_i32 s2, s2, s4
+; GFX7-GISEL-NEXT: buffer_store_dword v0, v0, s[12:15], 0 offen
+; GFX7-GISEL-NEXT: s_waitcnt vmcnt(0)
+; GFX7-GISEL-NEXT: v_mov_b32_e32 v0, s2
+; GFX7-GISEL-NEXT: s_mov_b32 s2, -1
+; GFX7-GISEL-NEXT: s_mov_b32 s3, 0xf000
+; GFX7-GISEL-NEXT: buffer_store_dword v0, off, s[0:3], 0
+; GFX7-GISEL-NEXT: s_endpgm
+;
+; GFX12-5-SDAG-LABEL: v_sad_u32_multi_use_sub_pat2:
+; GFX12-5-SDAG: ; %bb.0:
+; GFX12-5-SDAG-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-SDAG-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-SDAG-NEXT: v_nop
+; GFX12-5-SDAG-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-SDAG-NEXT: s_clause 0x1
+; GFX12-5-SDAG-NEXT: s_load_b96 s[0:2], s[4:5], 0x2c nv
+; GFX12-5-SDAG-NEXT: s_load_b64 s[6:7], s[4:5], 0x24 nv
+; GFX12-5-SDAG-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-SDAG-NEXT: s_min_u32 s3, s1, s0
+; GFX12-5-SDAG-NEXT: s_max_u32 s4, s1, s0
+; GFX12-5-SDAG-NEXT: s_sub_co_i32 s0, s0, s1
+; GFX12-5-SDAG-NEXT: s_sub_co_i32 s3, s4, s3
+; GFX12-5-SDAG-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s0
+; GFX12-5-SDAG-NEXT: s_add_co_i32 s1, s3, s2
+; GFX12-5-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
+; GFX12-5-SDAG-NEXT: v_mov_b32_e32 v2, s1
+; GFX12-5-SDAG-NEXT: scratch_store_b32 off, v1, s0 scope:SCOPE_SYS
+; GFX12-5-SDAG-NEXT: s_wait_storecnt 0x0
+; GFX12-5-SDAG-NEXT: global_store_b32 v0, v2, s[6:7]
+; GFX12-5-SDAG-NEXT: s_endpgm
+;
+; GFX12-5-GISEL-LABEL: v_sad_u32_multi_use_sub_pat2:
+; GFX12-5-GISEL: ; %bb.0:
+; GFX12-5-GISEL-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-GISEL-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-GISEL-NEXT: v_nop
+; GFX12-5-GISEL-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-GISEL-NEXT: s_clause 0x1
+; GFX12-5-GISEL-NEXT: s_load_b96 s[0:2], s[4:5], 0x2c nv
+; GFX12-5-GISEL-NEXT: s_load_b64 s[6:7], s[4:5], 0x24 nv
+; GFX12-5-GISEL-NEXT: v_mov_b32_e32 v2, 0
+; GFX12-5-GISEL-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-GISEL-NEXT: s_sub_co_u32 s3, s1, s0
+; GFX12-5-GISEL-NEXT: s_cselect_b32 s4, 1, 0
+; GFX12-5-GISEL-NEXT: s_sub_co_i32 s0, s0, s1
+; GFX12-5-GISEL-NEXT: s_cmp_lg_u32 s4, 0
+; GFX12-5-GISEL-NEXT: s_cselect_b32 s1, s0, s3
+; GFX12-5-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
+; GFX12-5-GISEL-NEXT: s_add_co_i32 s1, s1, s2
+; GFX12-5-GISEL-NEXT: v_dual_mov_b32 v0, s0 :: v_dual_mov_b32 v1, s1
+; GFX12-5-GISEL-NEXT: scratch_store_b32 off, v0, s0 scope:SCOPE_SYS
+; GFX12-5-GISEL-NEXT: s_wait_storecnt 0x0
+; GFX12-5-GISEL-NEXT: global_store_b32 v2, v1, s[6:7]
+; GFX12-5-GISEL-NEXT: s_endpgm
%icmp0 = icmp ugt i32 %a, %b
%sub0 = sub i32 %a, %b
store volatile i32 %sub0, ptr addrspace(5) poison
@@ -273,30 +779,102 @@ define amdgpu_kernel void @v_sad_u32_multi_use_sub_pat2(ptr addrspace(1) %out, i
}
define amdgpu_kernel void @v_sad_u32_multi_use_select_pat2(ptr addrspace(1) %out, i32 %a, i32 %b, i32 %c) {
-; GCN-LABEL: v_sad_u32_multi_use_select_pat2:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_mov_b64 s[22:23], s[2:3]
-; GCN-NEXT: s_mov_b64 s[20:21], s[0:1]
-; GCN-NEXT: s_load_dwordx4 s[0:3], s[8:9], 0x2
-; GCN-NEXT: s_load_dwordx2 s[4:5], s[8:9], 0x0
-; GCN-NEXT: s_add_i32 s12, s12, s17
-; GCN-NEXT: s_lshr_b32 flat_scratch_hi, s12, 8
-; GCN-NEXT: s_add_u32 s20, s20, s17
-; GCN-NEXT: s_addc_u32 s21, s21, 0
-; GCN-NEXT: s_waitcnt lgkmcnt(0)
-; GCN-NEXT: s_min_u32 s3, s1, s0
-; GCN-NEXT: s_max_u32 s0, s1, s0
-; GCN-NEXT: s_sub_i32 s0, s0, s3
-; GCN-NEXT: v_mov_b32_e32 v0, s4
-; GCN-NEXT: v_mov_b32_e32 v2, s0
-; GCN-NEXT: s_add_i32 s0, s0, s2
-; GCN-NEXT: s_mov_b32 flat_scratch_lo, s13
-; GCN-NEXT: v_mov_b32_e32 v1, s5
-; GCN-NEXT: buffer_store_dword v2, v0, s[20:23], 0 offen
-; GCN-NEXT: s_waitcnt vmcnt(0)
-; GCN-NEXT: v_mov_b32_e32 v2, s0
-; GCN-NEXT: flat_store_dword v[0:1], v2
-; GCN-NEXT: s_endpgm
+; GFX7-SDAG-LABEL: v_sad_u32_multi_use_select_pat2:
+; GFX7-SDAG: ; %bb.0:
+; GFX7-SDAG-NEXT: s_mov_b32 s12, SCRATCH_RSRC_DWORD0
+; GFX7-SDAG-NEXT: s_mov_b32 s13, SCRATCH_RSRC_DWORD1
+; GFX7-SDAG-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0xb
+; GFX7-SDAG-NEXT: s_load_dwordx2 s[4:5], s[4:5], 0x9
+; GFX7-SDAG-NEXT: s_mov_b32 s14, -1
+; GFX7-SDAG-NEXT: s_mov_b32 s15, 0xe8f000
+; GFX7-SDAG-NEXT: s_add_u32 s12, s12, s11
+; GFX7-SDAG-NEXT: s_addc_u32 s13, s13, 0
+; GFX7-SDAG-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-SDAG-NEXT: s_min_u32 s3, s1, s0
+; GFX7-SDAG-NEXT: s_max_u32 s0, s1, s0
+; GFX7-SDAG-NEXT: s_sub_i32 s0, s0, s3
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v0, s0
+; GFX7-SDAG-NEXT: s_add_i32 s0, s0, s2
+; GFX7-SDAG-NEXT: s_mov_b32 s7, 0xf000
+; GFX7-SDAG-NEXT: s_mov_b32 s6, -1
+; GFX7-SDAG-NEXT: buffer_store_dword v0, v0, s[12:15], 0 offen
+; GFX7-SDAG-NEXT: s_waitcnt vmcnt(0)
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v0, s0
+; GFX7-SDAG-NEXT: buffer_store_dword v0, off, s[4:7], 0
+; GFX7-SDAG-NEXT: s_endpgm
+;
+; GFX7-GISEL-LABEL: v_sad_u32_multi_use_select_pat2:
+; GFX7-GISEL: ; %bb.0:
+; GFX7-GISEL-NEXT: s_mov_b32 s12, SCRATCH_RSRC_DWORD0
+; GFX7-GISEL-NEXT: s_mov_b32 s13, SCRATCH_RSRC_DWORD1
+; GFX7-GISEL-NEXT: s_mov_b32 s14, -1
+; GFX7-GISEL-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x9
+; GFX7-GISEL-NEXT: s_load_dword s4, s[4:5], 0xd
+; GFX7-GISEL-NEXT: s_mov_b32 s15, 0xe8f000
+; GFX7-GISEL-NEXT: s_add_u32 s12, s12, s11
+; GFX7-GISEL-NEXT: s_addc_u32 s13, s13, 0
+; GFX7-GISEL-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-GISEL-NEXT: s_sub_u32 s5, s3, s2
+; GFX7-GISEL-NEXT: s_cselect_b32 s6, 1, 0
+; GFX7-GISEL-NEXT: s_sub_i32 s2, s2, s3
+; GFX7-GISEL-NEXT: s_cmp_lg_u32 s6, 0
+; GFX7-GISEL-NEXT: s_cselect_b32 s2, s2, s5
+; GFX7-GISEL-NEXT: v_mov_b32_e32 v0, s2
+; GFX7-GISEL-NEXT: s_add_i32 s2, s2, s4
+; GFX7-GISEL-NEXT: buffer_store_dword v0, v0, s[12:15], 0 offen
+; GFX7-GISEL-NEXT: s_waitcnt vmcnt(0)
+; GFX7-GISEL-NEXT: v_mov_b32_e32 v0, s2
+; GFX7-GISEL-NEXT: s_mov_b32 s2, -1
+; GFX7-GISEL-NEXT: s_mov_b32 s3, 0xf000
+; GFX7-GISEL-NEXT: buffer_store_dword v0, off, s[0:3], 0
+; GFX7-GISEL-NEXT: s_endpgm
+;
+; GFX12-5-SDAG-LABEL: v_sad_u32_multi_use_select_pat2:
+; GFX12-5-SDAG: ; %bb.0:
+; GFX12-5-SDAG-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-SDAG-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-SDAG-NEXT: v_nop
+; GFX12-5-SDAG-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-SDAG-NEXT: s_clause 0x1
+; GFX12-5-SDAG-NEXT: s_load_b96 s[0:2], s[4:5], 0x2c nv
+; GFX12-5-SDAG-NEXT: s_load_b64 s[6:7], s[4:5], 0x24 nv
+; GFX12-5-SDAG-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-SDAG-NEXT: s_min_u32 s3, s1, s0
+; GFX12-5-SDAG-NEXT: s_max_u32 s0, s1, s0
+; GFX12-5-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
+; GFX12-5-SDAG-NEXT: s_sub_co_i32 s0, s0, s3
+; GFX12-5-SDAG-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s0
+; GFX12-5-SDAG-NEXT: s_add_co_i32 s1, s0, s2
+; GFX12-5-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
+; GFX12-5-SDAG-NEXT: v_mov_b32_e32 v2, s1
+; GFX12-5-SDAG-NEXT: scratch_store_b32 off, v1, s0 scope:SCOPE_SYS
+; GFX12-5-SDAG-NEXT: s_wait_storecnt 0x0
+; GFX12-5-SDAG-NEXT: global_store_b32 v0, v2, s[6:7]
+; GFX12-5-SDAG-NEXT: s_endpgm
+;
+; GFX12-5-GISEL-LABEL: v_sad_u32_multi_use_select_pat2:
+; GFX12-5-GISEL: ; %bb.0:
+; GFX12-5-GISEL-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-GISEL-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-GISEL-NEXT: v_nop
+; GFX12-5-GISEL-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-GISEL-NEXT: s_clause 0x1
+; GFX12-5-GISEL-NEXT: s_load_b96 s[0:2], s[4:5], 0x2c nv
+; GFX12-5-GISEL-NEXT: s_load_b64 s[6:7], s[4:5], 0x24 nv
+; GFX12-5-GISEL-NEXT: v_mov_b32_e32 v2, 0
+; GFX12-5-GISEL-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-GISEL-NEXT: s_sub_co_u32 s3, s1, s0
+; GFX12-5-GISEL-NEXT: s_cselect_b32 s4, 1, 0
+; GFX12-5-GISEL-NEXT: s_sub_co_i32 s0, s0, s1
+; GFX12-5-GISEL-NEXT: s_cmp_lg_u32 s4, 0
+; GFX12-5-GISEL-NEXT: s_cselect_b32 s0, s0, s3
+; GFX12-5-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
+; GFX12-5-GISEL-NEXT: s_add_co_i32 s1, s0, s2
+; GFX12-5-GISEL-NEXT: v_dual_mov_b32 v0, s0 :: v_dual_mov_b32 v1, s1
+; GFX12-5-GISEL-NEXT: scratch_store_b32 off, v0, s0 scope:SCOPE_SYS
+; GFX12-5-GISEL-NEXT: s_wait_storecnt 0x0
+; GFX12-5-GISEL-NEXT: global_store_b32 v2, v1, s[6:7]
+; GFX12-5-GISEL-NEXT: s_endpgm
%icmp0 = icmp ugt i32 %a, %b
%sub0 = sub i32 %a, %b
%sub1 = sub i32 %b, %a
@@ -310,31 +888,125 @@ define amdgpu_kernel void @v_sad_u32_multi_use_select_pat2(ptr addrspace(1) %out
}
define amdgpu_kernel void @v_sad_u32_vector_pat1(ptr addrspace(1) %out, <4 x i32> %a, <4 x i32> %b, <4 x i32> %c) {
-; GCN-LABEL: v_sad_u32_vector_pat1:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_add_i32 s12, s12, s17
-; GCN-NEXT: s_mov_b32 flat_scratch_lo, s13
-; GCN-NEXT: s_lshr_b32 flat_scratch_hi, s12, 8
-; GCN-NEXT: s_load_dwordx8 s[0:7], s[8:9], 0x4
-; GCN-NEXT: s_load_dwordx4 s[12:15], s[8:9], 0xc
-; GCN-NEXT: s_load_dwordx2 s[8:9], s[8:9], 0x0
-; GCN-NEXT: s_waitcnt lgkmcnt(0)
-; GCN-NEXT: v_mov_b32_e32 v0, s7
-; GCN-NEXT: v_mov_b32_e32 v1, s15
-; GCN-NEXT: v_mov_b32_e32 v2, s6
-; GCN-NEXT: v_sad_u32 v3, s3, v0, v1
-; GCN-NEXT: v_mov_b32_e32 v0, s14
-; GCN-NEXT: v_sad_u32 v2, s2, v2, v0
-; GCN-NEXT: v_mov_b32_e32 v0, s5
-; GCN-NEXT: v_mov_b32_e32 v1, s13
-; GCN-NEXT: v_sad_u32 v1, s1, v0, v1
-; GCN-NEXT: v_mov_b32_e32 v0, s4
-; GCN-NEXT: v_mov_b32_e32 v4, s12
-; GCN-NEXT: v_sad_u32 v0, s0, v0, v4
-; GCN-NEXT: v_mov_b32_e32 v4, s8
-; GCN-NEXT: v_mov_b32_e32 v5, s9
-; GCN-NEXT: flat_store_dwordx4 v[4:5], v[0:3]
-; GCN-NEXT: s_endpgm
+; GFX7-SDAG-LABEL: v_sad_u32_vector_pat1:
+; GFX7-SDAG: ; %bb.0:
+; GFX7-SDAG-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0xd
+; GFX7-SDAG-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x15
+; GFX7-SDAG-NEXT: s_load_dwordx2 s[4:5], s[4:5], 0x9
+; GFX7-SDAG-NEXT: s_mov_b32 s7, 0xf000
+; GFX7-SDAG-NEXT: s_mov_b32 s6, -1
+; GFX7-SDAG-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v0, s15
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v1, s3
+; GFX7-SDAG-NEXT: v_sad_u32 v3, s11, v0, v1
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v0, s14
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v1, s2
+; GFX7-SDAG-NEXT: v_sad_u32 v2, s10, v0, v1
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v0, s13
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v1, s1
+; GFX7-SDAG-NEXT: v_sad_u32 v1, s9, v0, v1
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v0, s12
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v4, s0
+; GFX7-SDAG-NEXT: v_sad_u32 v0, s8, v0, v4
+; GFX7-SDAG-NEXT: buffer_store_dwordx4 v[0:3], off, s[4:7], 0
+; GFX7-SDAG-NEXT: s_endpgm
+;
+; GFX7-GISEL-LABEL: v_sad_u32_vector_pat1:
+; GFX7-GISEL: ; %bb.0:
+; GFX7-GISEL-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0xd
+; GFX7-GISEL-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x9
+; GFX7-GISEL-NEXT: s_load_dwordx4 s[4:7], s[4:5], 0x15
+; GFX7-GISEL-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-GISEL-NEXT: s_max_u32 s2, s8, s12
+; GFX7-GISEL-NEXT: s_max_u32 s3, s9, s13
+; GFX7-GISEL-NEXT: s_max_u32 s16, s10, s14
+; GFX7-GISEL-NEXT: s_max_u32 s17, s11, s15
+; GFX7-GISEL-NEXT: s_min_u32 s8, s8, s12
+; GFX7-GISEL-NEXT: s_min_u32 s9, s9, s13
+; GFX7-GISEL-NEXT: s_min_u32 s10, s10, s14
+; GFX7-GISEL-NEXT: s_min_u32 s11, s11, s15
+; GFX7-GISEL-NEXT: s_sub_i32 s2, s2, s8
+; GFX7-GISEL-NEXT: s_sub_i32 s3, s3, s9
+; GFX7-GISEL-NEXT: s_sub_i32 s8, s16, s10
+; GFX7-GISEL-NEXT: s_sub_i32 s9, s17, s11
+; GFX7-GISEL-NEXT: s_add_i32 s4, s2, s4
+; GFX7-GISEL-NEXT: s_add_i32 s5, s3, s5
+; GFX7-GISEL-NEXT: s_add_i32 s6, s8, s6
+; GFX7-GISEL-NEXT: s_add_i32 s7, s9, s7
+; GFX7-GISEL-NEXT: v_mov_b32_e32 v0, s4
+; GFX7-GISEL-NEXT: v_mov_b32_e32 v1, s5
+; GFX7-GISEL-NEXT: v_mov_b32_e32 v2, s6
+; GFX7-GISEL-NEXT: v_mov_b32_e32 v3, s7
+; GFX7-GISEL-NEXT: s_mov_b32 s2, -1
+; GFX7-GISEL-NEXT: s_mov_b32 s3, 0xf000
+; GFX7-GISEL-NEXT: buffer_store_dwordx4 v[0:3], off, s[0:3], 0
+; GFX7-GISEL-NEXT: s_endpgm
+;
+; GFX12-5-SDAG-LABEL: v_sad_u32_vector_pat1:
+; GFX12-5-SDAG: ; %bb.0:
+; GFX12-5-SDAG-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-SDAG-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-SDAG-NEXT: v_nop
+; GFX12-5-SDAG-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-SDAG-NEXT: s_clause 0x2
+; GFX12-5-SDAG-NEXT: s_load_b256 s[8:15], s[4:5], 0x34 nv
+; GFX12-5-SDAG-NEXT: s_load_b128 s[0:3], s[4:5], 0x54 nv
+; GFX12-5-SDAG-NEXT: s_load_b64 s[6:7], s[4:5], 0x24 nv
+; GFX12-5-SDAG-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-SDAG-NEXT: s_min_u32 s4, s8, s12
+; GFX12-5-SDAG-NEXT: s_max_u32 s5, s8, s12
+; GFX12-5-SDAG-NEXT: s_min_u32 s8, s9, s13
+; GFX12-5-SDAG-NEXT: s_max_u32 s9, s9, s13
+; GFX12-5-SDAG-NEXT: s_min_u32 s12, s10, s14
+; GFX12-5-SDAG-NEXT: s_max_u32 s10, s10, s14
+; GFX12-5-SDAG-NEXT: s_min_u32 s13, s11, s15
+; GFX12-5-SDAG-NEXT: s_max_u32 s11, s11, s15
+; GFX12-5-SDAG-NEXT: s_sub_co_i32 s4, s5, s4
+; GFX12-5-SDAG-NEXT: s_sub_co_i32 s5, s9, s8
+; GFX12-5-SDAG-NEXT: s_sub_co_i32 s8, s11, s13
+; GFX12-5-SDAG-NEXT: s_sub_co_i32 s9, s10, s12
+; GFX12-5-SDAG-NEXT: s_add_co_i32 s8, s8, s3
+; GFX12-5-SDAG-NEXT: s_add_co_i32 s9, s9, s2
+; GFX12-5-SDAG-NEXT: s_add_co_i32 s4, s4, s0
+; GFX12-5-SDAG-NEXT: s_add_co_i32 s5, s5, s1
+; GFX12-5-SDAG-NEXT: v_dual_mov_b32 v4, 0 :: v_dual_mov_b32 v0, s4
+; GFX12-5-SDAG-NEXT: v_dual_mov_b32 v1, s5 :: v_dual_mov_b32 v2, s9
+; GFX12-5-SDAG-NEXT: v_mov_b32_e32 v3, s8
+; GFX12-5-SDAG-NEXT: global_store_b128 v4, v[0:3], s[6:7]
+; GFX12-5-SDAG-NEXT: s_endpgm
+;
+; GFX12-5-GISEL-LABEL: v_sad_u32_vector_pat1:
+; GFX12-5-GISEL: ; %bb.0:
+; GFX12-5-GISEL-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-GISEL-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-GISEL-NEXT: v_nop
+; GFX12-5-GISEL-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-GISEL-NEXT: s_clause 0x2
+; GFX12-5-GISEL-NEXT: s_load_b256 s[8:15], s[4:5], 0x34 nv
+; GFX12-5-GISEL-NEXT: s_load_b128 s[0:3], s[4:5], 0x54 nv
+; GFX12-5-GISEL-NEXT: s_load_b64 s[6:7], s[4:5], 0x24 nv
+; GFX12-5-GISEL-NEXT: v_mov_b32_e32 v4, 0
+; GFX12-5-GISEL-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-GISEL-NEXT: s_max_u32 s4, s8, s12
+; GFX12-5-GISEL-NEXT: s_max_u32 s5, s9, s13
+; GFX12-5-GISEL-NEXT: s_max_u32 s16, s10, s14
+; GFX12-5-GISEL-NEXT: s_max_u32 s17, s11, s15
+; GFX12-5-GISEL-NEXT: s_min_u32 s8, s8, s12
+; GFX12-5-GISEL-NEXT: s_min_u32 s9, s9, s13
+; GFX12-5-GISEL-NEXT: s_min_u32 s10, s10, s14
+; GFX12-5-GISEL-NEXT: s_min_u32 s11, s11, s15
+; GFX12-5-GISEL-NEXT: s_sub_co_i32 s4, s4, s8
+; GFX12-5-GISEL-NEXT: s_sub_co_i32 s5, s5, s9
+; GFX12-5-GISEL-NEXT: s_sub_co_i32 s8, s16, s10
+; GFX12-5-GISEL-NEXT: s_sub_co_i32 s9, s17, s11
+; GFX12-5-GISEL-NEXT: s_add_co_i32 s0, s4, s0
+; GFX12-5-GISEL-NEXT: s_add_co_i32 s1, s5, s1
+; GFX12-5-GISEL-NEXT: s_add_co_i32 s2, s8, s2
+; GFX12-5-GISEL-NEXT: s_add_co_i32 s3, s9, s3
+; GFX12-5-GISEL-NEXT: v_mov_b64_e32 v[0:1], s[0:1]
+; GFX12-5-GISEL-NEXT: v_mov_b64_e32 v[2:3], s[2:3]
+; GFX12-5-GISEL-NEXT: global_store_b128 v4, v[0:3], s[6:7]
+; GFX12-5-GISEL-NEXT: s_endpgm
%icmp0 = icmp ugt <4 x i32> %a, %b
%t0 = select <4 x i1> %icmp0, <4 x i32> %a, <4 x i32> %b
@@ -349,31 +1021,149 @@ define amdgpu_kernel void @v_sad_u32_vector_pat1(ptr addrspace(1) %out, <4 x i32
}
define amdgpu_kernel void @v_sad_u32_vector_pat2(ptr addrspace(1) %out, <4 x i32> %a, <4 x i32> %b, <4 x i32> %c) {
-; GCN-LABEL: v_sad_u32_vector_pat2:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_add_i32 s12, s12, s17
-; GCN-NEXT: s_mov_b32 flat_scratch_lo, s13
-; GCN-NEXT: s_lshr_b32 flat_scratch_hi, s12, 8
-; GCN-NEXT: s_load_dwordx8 s[0:7], s[8:9], 0x4
-; GCN-NEXT: s_load_dwordx4 s[12:15], s[8:9], 0xc
-; GCN-NEXT: s_load_dwordx2 s[8:9], s[8:9], 0x0
-; GCN-NEXT: s_waitcnt lgkmcnt(0)
-; GCN-NEXT: v_mov_b32_e32 v0, s7
-; GCN-NEXT: v_mov_b32_e32 v1, s15
-; GCN-NEXT: v_mov_b32_e32 v2, s6
-; GCN-NEXT: v_sad_u32 v3, s3, v0, v1
-; GCN-NEXT: v_mov_b32_e32 v0, s14
-; GCN-NEXT: v_sad_u32 v2, s2, v2, v0
-; GCN-NEXT: v_mov_b32_e32 v0, s5
-; GCN-NEXT: v_mov_b32_e32 v1, s13
-; GCN-NEXT: v_sad_u32 v1, s1, v0, v1
-; GCN-NEXT: v_mov_b32_e32 v0, s4
-; GCN-NEXT: v_mov_b32_e32 v4, s12
-; GCN-NEXT: v_sad_u32 v0, s0, v0, v4
-; GCN-NEXT: v_mov_b32_e32 v4, s8
-; GCN-NEXT: v_mov_b32_e32 v5, s9
-; GCN-NEXT: flat_store_dwordx4 v[4:5], v[0:3]
-; GCN-NEXT: s_endpgm
+; GFX7-SDAG-LABEL: v_sad_u32_vector_pat2:
+; GFX7-SDAG: ; %bb.0:
+; GFX7-SDAG-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0xd
+; GFX7-SDAG-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x15
+; GFX7-SDAG-NEXT: s_load_dwordx2 s[4:5], s[4:5], 0x9
+; GFX7-SDAG-NEXT: s_mov_b32 s7, 0xf000
+; GFX7-SDAG-NEXT: s_mov_b32 s6, -1
+; GFX7-SDAG-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v0, s15
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v1, s3
+; GFX7-SDAG-NEXT: v_sad_u32 v3, s11, v0, v1
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v0, s14
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v1, s2
+; GFX7-SDAG-NEXT: v_sad_u32 v2, s10, v0, v1
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v0, s13
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v1, s1
+; GFX7-SDAG-NEXT: v_sad_u32 v1, s9, v0, v1
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v0, s12
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v4, s0
+; GFX7-SDAG-NEXT: v_sad_u32 v0, s8, v0, v4
+; GFX7-SDAG-NEXT: buffer_store_dwordx4 v[0:3], off, s[4:7], 0
+; GFX7-SDAG-NEXT: s_endpgm
+;
+; GFX7-GISEL-LABEL: v_sad_u32_vector_pat2:
+; GFX7-GISEL: ; %bb.0:
+; GFX7-GISEL-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0xd
+; GFX7-GISEL-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x9
+; GFX7-GISEL-NEXT: s_load_dwordx4 s[4:7], s[4:5], 0x15
+; GFX7-GISEL-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-GISEL-NEXT: s_cmp_gt_u32 s8, s12
+; GFX7-GISEL-NEXT: s_cselect_b32 s2, 1, 0
+; GFX7-GISEL-NEXT: s_cmp_gt_u32 s9, s13
+; GFX7-GISEL-NEXT: s_cselect_b32 s3, 1, 0
+; GFX7-GISEL-NEXT: s_cmp_gt_u32 s10, s14
+; GFX7-GISEL-NEXT: s_cselect_b32 s16, 1, 0
+; GFX7-GISEL-NEXT: s_cmp_gt_u32 s11, s15
+; GFX7-GISEL-NEXT: s_cselect_b32 s17, 1, 0
+; GFX7-GISEL-NEXT: s_sub_i32 s18, s8, s12
+; GFX7-GISEL-NEXT: s_sub_i32 s19, s9, s13
+; GFX7-GISEL-NEXT: s_sub_i32 s20, s10, s14
+; GFX7-GISEL-NEXT: s_sub_i32 s21, s11, s15
+; GFX7-GISEL-NEXT: s_sub_i32 s8, s12, s8
+; GFX7-GISEL-NEXT: s_sub_i32 s9, s13, s9
+; GFX7-GISEL-NEXT: s_sub_i32 s10, s14, s10
+; GFX7-GISEL-NEXT: s_sub_i32 s11, s15, s11
+; GFX7-GISEL-NEXT: s_cmp_lg_u32 s2, 0
+; GFX7-GISEL-NEXT: s_cselect_b32 s2, s18, s8
+; GFX7-GISEL-NEXT: s_cmp_lg_u32 s3, 0
+; GFX7-GISEL-NEXT: s_cselect_b32 s3, s19, s9
+; GFX7-GISEL-NEXT: s_cmp_lg_u32 s16, 0
+; GFX7-GISEL-NEXT: s_cselect_b32 s8, s20, s10
+; GFX7-GISEL-NEXT: s_cmp_lg_u32 s17, 0
+; GFX7-GISEL-NEXT: s_cselect_b32 s9, s21, s11
+; GFX7-GISEL-NEXT: s_add_i32 s4, s2, s4
+; GFX7-GISEL-NEXT: s_add_i32 s5, s3, s5
+; GFX7-GISEL-NEXT: s_add_i32 s6, s8, s6
+; GFX7-GISEL-NEXT: s_add_i32 s7, s9, s7
+; GFX7-GISEL-NEXT: v_mov_b32_e32 v0, s4
+; GFX7-GISEL-NEXT: v_mov_b32_e32 v1, s5
+; GFX7-GISEL-NEXT: v_mov_b32_e32 v2, s6
+; GFX7-GISEL-NEXT: v_mov_b32_e32 v3, s7
+; GFX7-GISEL-NEXT: s_mov_b32 s2, -1
+; GFX7-GISEL-NEXT: s_mov_b32 s3, 0xf000
+; GFX7-GISEL-NEXT: buffer_store_dwordx4 v[0:3], off, s[0:3], 0
+; GFX7-GISEL-NEXT: s_endpgm
+;
+; GFX12-5-SDAG-LABEL: v_sad_u32_vector_pat2:
+; GFX12-5-SDAG: ; %bb.0:
+; GFX12-5-SDAG-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-SDAG-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-SDAG-NEXT: v_nop
+; GFX12-5-SDAG-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-SDAG-NEXT: s_clause 0x2
+; GFX12-5-SDAG-NEXT: s_load_b256 s[8:15], s[4:5], 0x34 nv
+; GFX12-5-SDAG-NEXT: s_load_b128 s[0:3], s[4:5], 0x54 nv
+; GFX12-5-SDAG-NEXT: s_load_b64 s[6:7], s[4:5], 0x24 nv
+; GFX12-5-SDAG-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-SDAG-NEXT: s_min_u32 s4, s8, s12
+; GFX12-5-SDAG-NEXT: s_max_u32 s5, s8, s12
+; GFX12-5-SDAG-NEXT: s_min_u32 s8, s9, s13
+; GFX12-5-SDAG-NEXT: s_max_u32 s9, s9, s13
+; GFX12-5-SDAG-NEXT: s_min_u32 s12, s10, s14
+; GFX12-5-SDAG-NEXT: s_max_u32 s10, s10, s14
+; GFX12-5-SDAG-NEXT: s_min_u32 s13, s11, s15
+; GFX12-5-SDAG-NEXT: s_max_u32 s11, s11, s15
+; GFX12-5-SDAG-NEXT: s_sub_co_i32 s4, s5, s4
+; GFX12-5-SDAG-NEXT: s_sub_co_i32 s5, s9, s8
+; GFX12-5-SDAG-NEXT: s_sub_co_i32 s8, s11, s13
+; GFX12-5-SDAG-NEXT: s_sub_co_i32 s9, s10, s12
+; GFX12-5-SDAG-NEXT: s_add_co_i32 s8, s8, s3
+; GFX12-5-SDAG-NEXT: s_add_co_i32 s9, s9, s2
+; GFX12-5-SDAG-NEXT: s_add_co_i32 s4, s4, s0
+; GFX12-5-SDAG-NEXT: s_add_co_i32 s5, s5, s1
+; GFX12-5-SDAG-NEXT: v_dual_mov_b32 v4, 0 :: v_dual_mov_b32 v0, s4
+; GFX12-5-SDAG-NEXT: v_dual_mov_b32 v1, s5 :: v_dual_mov_b32 v2, s9
+; GFX12-5-SDAG-NEXT: v_mov_b32_e32 v3, s8
+; GFX12-5-SDAG-NEXT: global_store_b128 v4, v[0:3], s[6:7]
+; GFX12-5-SDAG-NEXT: s_endpgm
+;
+; GFX12-5-GISEL-LABEL: v_sad_u32_vector_pat2:
+; GFX12-5-GISEL: ; %bb.0:
+; GFX12-5-GISEL-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-GISEL-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-GISEL-NEXT: v_nop
+; GFX12-5-GISEL-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-GISEL-NEXT: s_clause 0x2
+; GFX12-5-GISEL-NEXT: s_load_b256 s[8:15], s[4:5], 0x34 nv
+; GFX12-5-GISEL-NEXT: s_load_b128 s[0:3], s[4:5], 0x54 nv
+; GFX12-5-GISEL-NEXT: s_load_b64 s[6:7], s[4:5], 0x24 nv
+; GFX12-5-GISEL-NEXT: v_mov_b32_e32 v4, 0
+; GFX12-5-GISEL-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-GISEL-NEXT: s_cmp_gt_u32 s8, s12
+; GFX12-5-GISEL-NEXT: s_cselect_b32 s16, 1, 0
+; GFX12-5-GISEL-NEXT: s_cmp_gt_u32 s9, s13
+; GFX12-5-GISEL-NEXT: s_cselect_b32 s17, 1, 0
+; GFX12-5-GISEL-NEXT: s_cmp_gt_u32 s10, s14
+; GFX12-5-GISEL-NEXT: s_cselect_b32 s18, 1, 0
+; GFX12-5-GISEL-NEXT: s_cmp_gt_u32 s11, s15
+; GFX12-5-GISEL-NEXT: s_cselect_b32 s19, 1, 0
+; GFX12-5-GISEL-NEXT: s_sub_co_i32 s4, s8, s12
+; GFX12-5-GISEL-NEXT: s_sub_co_i32 s5, s9, s13
+; GFX12-5-GISEL-NEXT: s_sub_co_i32 s20, s10, s14
+; GFX12-5-GISEL-NEXT: s_sub_co_i32 s21, s11, s15
+; GFX12-5-GISEL-NEXT: s_sub_co_i32 s8, s12, s8
+; GFX12-5-GISEL-NEXT: s_sub_co_i32 s9, s13, s9
+; GFX12-5-GISEL-NEXT: s_sub_co_i32 s10, s14, s10
+; GFX12-5-GISEL-NEXT: s_sub_co_i32 s11, s15, s11
+; GFX12-5-GISEL-NEXT: s_cmp_lg_u32 s16, 0
+; GFX12-5-GISEL-NEXT: s_cselect_b32 s4, s4, s8
+; GFX12-5-GISEL-NEXT: s_cmp_lg_u32 s17, 0
+; GFX12-5-GISEL-NEXT: s_cselect_b32 s5, s5, s9
+; GFX12-5-GISEL-NEXT: s_cmp_lg_u32 s18, 0
+; GFX12-5-GISEL-NEXT: s_cselect_b32 s8, s20, s10
+; GFX12-5-GISEL-NEXT: s_cmp_lg_u32 s19, 0
+; GFX12-5-GISEL-NEXT: s_cselect_b32 s9, s21, s11
+; GFX12-5-GISEL-NEXT: s_add_co_i32 s0, s4, s0
+; GFX12-5-GISEL-NEXT: s_add_co_i32 s1, s5, s1
+; GFX12-5-GISEL-NEXT: s_add_co_i32 s2, s8, s2
+; GFX12-5-GISEL-NEXT: s_add_co_i32 s3, s9, s3
+; GFX12-5-GISEL-NEXT: v_mov_b64_e32 v[0:1], s[0:1]
+; GFX12-5-GISEL-NEXT: v_mov_b64_e32 v[2:3], s[2:3]
+; GFX12-5-GISEL-NEXT: global_store_b128 v4, v[0:3], s[6:7]
+; GFX12-5-GISEL-NEXT: s_endpgm
%icmp0 = icmp ugt <4 x i32> %a, %b
%sub0 = sub <4 x i32> %a, %b
%sub1 = sub <4 x i32> %b, %a
@@ -386,22 +1176,79 @@ define amdgpu_kernel void @v_sad_u32_vector_pat2(ptr addrspace(1) %out, <4 x i32
}
define amdgpu_kernel void @v_sad_u32_i16_pat1(ptr addrspace(1) %out, i16 %a, i16 %b, i16 %c) {
-; GCN-LABEL: v_sad_u32_i16_pat1:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_load_dwordx4 s[0:3], s[8:9], 0x0
-; GCN-NEXT: s_add_i32 s12, s12, s17
-; GCN-NEXT: s_mov_b32 flat_scratch_lo, s13
-; GCN-NEXT: s_lshr_b32 flat_scratch_hi, s12, 8
-; GCN-NEXT: s_waitcnt lgkmcnt(0)
-; GCN-NEXT: s_and_b32 s4, s2, 0xffff
-; GCN-NEXT: s_lshr_b32 s2, s2, 16
-; GCN-NEXT: v_mov_b32_e32 v0, s3
-; GCN-NEXT: v_mov_b32_e32 v1, s2
-; GCN-NEXT: v_sad_u32 v2, s4, v1, v0
-; GCN-NEXT: v_mov_b32_e32 v0, s0
-; GCN-NEXT: v_mov_b32_e32 v1, s1
-; GCN-NEXT: flat_store_short v[0:1], v2
-; GCN-NEXT: s_endpgm
+; GFX7-SDAG-LABEL: v_sad_u32_i16_pat1:
+; GFX7-SDAG: ; %bb.0:
+; GFX7-SDAG-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x9
+; GFX7-SDAG-NEXT: s_mov_b32 s7, 0xf000
+; GFX7-SDAG-NEXT: s_mov_b32 s6, -1
+; GFX7-SDAG-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-SDAG-NEXT: s_mov_b32 s5, s1
+; GFX7-SDAG-NEXT: s_lshr_b32 s1, s2, 16
+; GFX7-SDAG-NEXT: s_mov_b32 s4, s0
+; GFX7-SDAG-NEXT: s_and_b32 s0, s2, 0xffff
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v0, s1
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v1, s3
+; GFX7-SDAG-NEXT: v_sad_u32 v0, s0, v0, v1
+; GFX7-SDAG-NEXT: buffer_store_short v0, off, s[4:7], 0
+; GFX7-SDAG-NEXT: s_endpgm
+;
+; GFX7-GISEL-LABEL: v_sad_u32_i16_pat1:
+; GFX7-GISEL: ; %bb.0:
+; GFX7-GISEL-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x9
+; GFX7-GISEL-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-GISEL-NEXT: s_lshr_b32 s4, s2, 16
+; GFX7-GISEL-NEXT: s_and_b32 s5, s2, 0xffff
+; GFX7-GISEL-NEXT: s_max_u32 s6, s5, s4
+; GFX7-GISEL-NEXT: s_min_u32 s4, s5, s4
+; GFX7-GISEL-NEXT: s_sub_i32 s4, s6, s4
+; GFX7-GISEL-NEXT: s_add_i32 s3, s4, s3
+; GFX7-GISEL-NEXT: s_mov_b32 s2, -1
+; GFX7-GISEL-NEXT: v_mov_b32_e32 v0, s3
+; GFX7-GISEL-NEXT: s_mov_b32 s3, 0xf000
+; GFX7-GISEL-NEXT: buffer_store_short v0, off, s[0:3], 0
+; GFX7-GISEL-NEXT: s_endpgm
+;
+; GFX12-5-SDAG-LABEL: v_sad_u32_i16_pat1:
+; GFX12-5-SDAG: ; %bb.0:
+; GFX12-5-SDAG-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-SDAG-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-SDAG-NEXT: v_nop
+; GFX12-5-SDAG-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-SDAG-NEXT: s_load_b128 s[0:3], s[4:5], 0x24 nv
+; GFX12-5-SDAG-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-SDAG-NEXT: s_lshr_b32 s4, s2, 16
+; GFX12-5-SDAG-NEXT: s_and_b32 s2, s2, 0xffff
+; GFX12-5-SDAG-NEXT: s_and_b32 s4, s4, 0xffff
+; GFX12-5-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
+; GFX12-5-SDAG-NEXT: s_min_u32 s5, s2, s4
+; GFX12-5-SDAG-NEXT: s_max_u32 s2, s2, s4
+; GFX12-5-SDAG-NEXT: s_sub_co_i32 s2, s2, s5
+; GFX12-5-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
+; GFX12-5-SDAG-NEXT: s_add_co_i32 s2, s2, s3
+; GFX12-5-SDAG-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s2
+; GFX12-5-SDAG-NEXT: global_store_b16 v0, v1, s[0:1]
+; GFX12-5-SDAG-NEXT: s_endpgm
+;
+; GFX12-5-GISEL-LABEL: v_sad_u32_i16_pat1:
+; GFX12-5-GISEL: ; %bb.0:
+; GFX12-5-GISEL-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-GISEL-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-GISEL-NEXT: v_nop
+; GFX12-5-GISEL-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-GISEL-NEXT: s_load_b128 s[0:3], s[4:5], 0x24 nv
+; GFX12-5-GISEL-NEXT: v_mov_b32_e32 v1, 0
+; GFX12-5-GISEL-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-GISEL-NEXT: s_lshr_b32 s4, s2, 16
+; GFX12-5-GISEL-NEXT: s_and_b32 s2, 0xffff, s2
+; GFX12-5-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
+; GFX12-5-GISEL-NEXT: s_max_u32 s5, s2, s4
+; GFX12-5-GISEL-NEXT: s_min_u32 s2, s2, s4
+; GFX12-5-GISEL-NEXT: s_sub_co_i32 s2, s5, s2
+; GFX12-5-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
+; GFX12-5-GISEL-NEXT: s_add_co_i32 s2, s2, s3
+; GFX12-5-GISEL-NEXT: v_mov_b16_e32 v0.l, s2
+; GFX12-5-GISEL-NEXT: global_store_b16 v1, v0, s[0:1]
+; GFX12-5-GISEL-NEXT: s_endpgm
%icmp0 = icmp ugt i16 %a, %b
%t0 = select i1 %icmp0, i16 %a, i16 %b
@@ -416,24 +1263,111 @@ define amdgpu_kernel void @v_sad_u32_i16_pat1(ptr addrspace(1) %out, i16 %a, i16
}
define amdgpu_kernel void @v_sad_u32_i16_pat2(ptr addrspace(1) %out) {
-; GCN-LABEL: v_sad_u32_i16_pat2:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_add_i32 s12, s12, s17
-; GCN-NEXT: s_mov_b32 flat_scratch_lo, s13
-; GCN-NEXT: s_lshr_b32 flat_scratch_hi, s12, 8
-; GCN-NEXT: flat_load_ushort v0, v[0:1] glc
-; GCN-NEXT: s_waitcnt vmcnt(0)
-; GCN-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x0
-; GCN-NEXT: flat_load_ushort v1, v[0:1] glc
-; GCN-NEXT: s_waitcnt vmcnt(0)
-; GCN-NEXT: flat_load_ushort v2, v[0:1] glc
-; GCN-NEXT: s_waitcnt vmcnt(0)
-; GCN-NEXT: v_sad_u32 v2, v0, v1, v2
-; GCN-NEXT: s_waitcnt lgkmcnt(0)
-; GCN-NEXT: v_mov_b32_e32 v0, s0
-; GCN-NEXT: v_mov_b32_e32 v1, s1
-; GCN-NEXT: flat_store_short v[0:1], v2
-; GCN-NEXT: s_endpgm
+; GFX7-SDAG-LABEL: v_sad_u32_i16_pat2:
+; GFX7-SDAG: ; %bb.0:
+; GFX7-SDAG-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x9
+; GFX7-SDAG-NEXT: s_mov_b32 s3, 0xf000
+; GFX7-SDAG-NEXT: s_mov_b32 s2, -1
+; GFX7-SDAG-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-SDAG-NEXT: buffer_load_ushort v0, off, s[0:3], 0 glc
+; GFX7-SDAG-NEXT: s_waitcnt vmcnt(0)
+; GFX7-SDAG-NEXT: buffer_load_ushort v1, off, s[0:3], 0 glc
+; GFX7-SDAG-NEXT: s_waitcnt vmcnt(0)
+; GFX7-SDAG-NEXT: buffer_load_ushort v2, off, s[0:3], 0 glc
+; GFX7-SDAG-NEXT: s_waitcnt vmcnt(0)
+; GFX7-SDAG-NEXT: v_sad_u32 v0, v0, v1, v2
+; GFX7-SDAG-NEXT: buffer_store_short v0, off, s[0:3], 0
+; GFX7-SDAG-NEXT: s_endpgm
+;
+; GFX7-GISEL-LABEL: v_sad_u32_i16_pat2:
+; GFX7-GISEL: ; %bb.0:
+; GFX7-GISEL-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x9
+; GFX7-GISEL-NEXT: s_mov_b32 s2, -1
+; GFX7-GISEL-NEXT: s_mov_b32 s3, 0xf000
+; GFX7-GISEL-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-GISEL-NEXT: buffer_load_ushort v0, off, s[0:3], 0 glc
+; GFX7-GISEL-NEXT: s_waitcnt vmcnt(0)
+; GFX7-GISEL-NEXT: buffer_load_ushort v1, off, s[0:3], 0 glc
+; GFX7-GISEL-NEXT: s_waitcnt vmcnt(0)
+; GFX7-GISEL-NEXT: buffer_load_ushort v2, off, s[0:3], 0 glc
+; GFX7-GISEL-NEXT: s_waitcnt vmcnt(0)
+; GFX7-GISEL-NEXT: v_readfirstlane_b32 s4, v0
+; GFX7-GISEL-NEXT: v_readfirstlane_b32 s5, v1
+; GFX7-GISEL-NEXT: s_cmp_gt_u32 s4, s5
+; GFX7-GISEL-NEXT: s_cselect_b32 s7, 1, 0
+; GFX7-GISEL-NEXT: s_sub_i32 s8, s4, s5
+; GFX7-GISEL-NEXT: s_sub_i32 s4, s5, s4
+; GFX7-GISEL-NEXT: s_cmp_lg_u32 s7, 0
+; GFX7-GISEL-NEXT: v_readfirstlane_b32 s6, v2
+; GFX7-GISEL-NEXT: s_cselect_b32 s4, s8, s4
+; GFX7-GISEL-NEXT: s_add_i32 s4, s4, s6
+; GFX7-GISEL-NEXT: v_mov_b32_e32 v0, s4
+; GFX7-GISEL-NEXT: buffer_store_short v0, off, s[0:3], 0
+; GFX7-GISEL-NEXT: s_endpgm
+;
+; GFX12-5-SDAG-LABEL: v_sad_u32_i16_pat2:
+; GFX12-5-SDAG: ; %bb.0:
+; GFX12-5-SDAG-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-SDAG-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-SDAG-NEXT: v_nop
+; GFX12-5-SDAG-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-SDAG-NEXT: global_load_u16 v0, v[0:1], off scope:SCOPE_SYS
+; GFX12-5-SDAG-NEXT: s_wait_loadcnt 0x0
+; GFX12-5-SDAG-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv
+; GFX12-5-SDAG-NEXT: s_wait_xcnt 0x0
+; GFX12-5-SDAG-NEXT: global_load_u16 v1, v[0:1], off scope:SCOPE_SYS
+; GFX12-5-SDAG-NEXT: s_wait_loadcnt 0x0
+; GFX12-5-SDAG-NEXT: ; kill: killed $vgpr0_vgpr1
+; GFX12-5-SDAG-NEXT: ; kill: killed $vgpr0_vgpr1
+; GFX12-5-SDAG-NEXT: ; kill: killed $vgpr0_vgpr1
+; GFX12-5-SDAG-NEXT: global_load_u16 v2, v[0:1], off scope:SCOPE_SYS
+; GFX12-5-SDAG-NEXT: s_wait_loadcnt 0x0
+; GFX12-5-SDAG-NEXT: v_and_b32_e32 v0, 0xffff, v0
+; GFX12-5-SDAG-NEXT: v_and_b32_e32 v1, 0xffff, v1
+; GFX12-5-SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
+; GFX12-5-SDAG-NEXT: v_min_u32_e32 v3, v0, v1
+; GFX12-5-SDAG-NEXT: v_max_u32_e32 v0, v0, v1
+; GFX12-5-SDAG-NEXT: v_dual_mov_b32 v1, 0 :: v_dual_sub_nc_u32 v0, v0, v3
+; GFX12-5-SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX12-5-SDAG-NEXT: v_add_nc_u32_e32 v0, v0, v2
+; GFX12-5-SDAG-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-SDAG-NEXT: global_store_b16 v1, v0, s[0:1]
+; GFX12-5-SDAG-NEXT: s_endpgm
+;
+; GFX12-5-GISEL-LABEL: v_sad_u32_i16_pat2:
+; GFX12-5-GISEL: ; %bb.0:
+; GFX12-5-GISEL-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-GISEL-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-GISEL-NEXT: v_nop
+; GFX12-5-GISEL-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-GISEL-NEXT: global_load_u16 v0, v[0:1], off scope:SCOPE_SYS
+; GFX12-5-GISEL-NEXT: s_wait_loadcnt 0x0
+; GFX12-5-GISEL-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv
+; GFX12-5-GISEL-NEXT: s_wait_xcnt 0x0
+; GFX12-5-GISEL-NEXT: global_load_u16 v1, v[0:1], off scope:SCOPE_SYS
+; GFX12-5-GISEL-NEXT: s_wait_loadcnt 0x0
+; GFX12-5-GISEL-NEXT: v_readfirstlane_b32 s2, v0
+; GFX12-5-GISEL-NEXT: ; kill: killed $vgpr0_vgpr1
+; GFX12-5-GISEL-NEXT: s_and_b32 s4, 0xffff, s2
+; GFX12-5-GISEL-NEXT: global_load_u16 v2, v[0:1], off scope:SCOPE_SYS
+; GFX12-5-GISEL-NEXT: s_wait_loadcnt 0x0
+; GFX12-5-GISEL-NEXT: v_readfirstlane_b32 s3, v1
+; GFX12-5-GISEL-NEXT: v_mov_b32_e32 v1, 0
+; GFX12-5-GISEL-NEXT: s_and_b32 s5, 0xffff, s3
+; GFX12-5-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
+; GFX12-5-GISEL-NEXT: s_cmp_gt_u32 s4, s5
+; GFX12-5-GISEL-NEXT: s_cselect_b32 s4, 1, 0
+; GFX12-5-GISEL-NEXT: s_sub_co_i32 s6, s2, s3
+; GFX12-5-GISEL-NEXT: s_sub_co_i32 s2, s3, s2
+; GFX12-5-GISEL-NEXT: s_cmp_lg_u32 s4, 0
+; GFX12-5-GISEL-NEXT: s_cselect_b32 s2, s6, s2
+; GFX12-5-GISEL-NEXT: v_readfirstlane_b32 s5, v2
+; GFX12-5-GISEL-NEXT: s_add_co_i32 s2, s2, s5
+; GFX12-5-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
+; GFX12-5-GISEL-NEXT: v_mov_b16_e32 v0.l, s2
+; GFX12-5-GISEL-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-GISEL-NEXT: global_store_b16 v1, v0, s[0:1]
+; GFX12-5-GISEL-NEXT: s_endpgm
%a = load volatile i16, ptr addrspace(1) poison
%b = load volatile i16, ptr addrspace(1) poison
%c = load volatile i16, ptr addrspace(1) poison
@@ -449,24 +1383,84 @@ define amdgpu_kernel void @v_sad_u32_i16_pat2(ptr addrspace(1) %out) {
}
define amdgpu_kernel void @v_sad_u32_i8_pat1(ptr addrspace(1) %out, i8 %a, i8 %b, i8 %c) {
-; GCN-LABEL: v_sad_u32_i8_pat1:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_load_dword s2, s[8:9], 0x2
-; GCN-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x0
-; GCN-NEXT: s_add_i32 s12, s12, s17
-; GCN-NEXT: s_mov_b32 flat_scratch_lo, s13
-; GCN-NEXT: s_lshr_b32 flat_scratch_hi, s12, 8
-; GCN-NEXT: s_waitcnt lgkmcnt(0)
-; GCN-NEXT: s_and_b32 s3, s2, 0xff
-; GCN-NEXT: s_bfe_u32 s4, s2, 0x80008
-; GCN-NEXT: s_lshr_b32 s2, s2, 16
-; GCN-NEXT: v_mov_b32_e32 v0, s4
-; GCN-NEXT: v_mov_b32_e32 v1, s2
-; GCN-NEXT: v_sad_u32 v2, s3, v0, v1
-; GCN-NEXT: v_mov_b32_e32 v0, s0
-; GCN-NEXT: v_mov_b32_e32 v1, s1
-; GCN-NEXT: flat_store_byte v[0:1], v2
-; GCN-NEXT: s_endpgm
+; GFX7-SDAG-LABEL: v_sad_u32_i8_pat1:
+; GFX7-SDAG: ; %bb.0:
+; GFX7-SDAG-NEXT: s_load_dword s6, s[4:5], 0xb
+; GFX7-SDAG-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x9
+; GFX7-SDAG-NEXT: s_mov_b32 s3, 0xf000
+; GFX7-SDAG-NEXT: s_mov_b32 s2, -1
+; GFX7-SDAG-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-SDAG-NEXT: s_and_b32 s4, s6, 0xff
+; GFX7-SDAG-NEXT: s_bfe_u32 s5, s6, 0x80008
+; GFX7-SDAG-NEXT: s_lshr_b32 s6, s6, 16
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v0, s5
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v1, s6
+; GFX7-SDAG-NEXT: v_sad_u32 v0, s4, v0, v1
+; GFX7-SDAG-NEXT: buffer_store_byte v0, off, s[0:3], 0
+; GFX7-SDAG-NEXT: s_endpgm
+;
+; GFX7-GISEL-LABEL: v_sad_u32_i8_pat1:
+; GFX7-GISEL: ; %bb.0:
+; GFX7-GISEL-NEXT: s_load_dword s2, s[4:5], 0xb
+; GFX7-GISEL-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x9
+; GFX7-GISEL-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-GISEL-NEXT: s_lshr_b32 s3, s2, 16
+; GFX7-GISEL-NEXT: s_and_b32 s4, s2, 0xff
+; GFX7-GISEL-NEXT: s_bfe_u32 s2, s2, 0x80008
+; GFX7-GISEL-NEXT: s_max_u32 s5, s4, s2
+; GFX7-GISEL-NEXT: s_min_u32 s2, s4, s2
+; GFX7-GISEL-NEXT: s_sub_i32 s2, s5, s2
+; GFX7-GISEL-NEXT: s_add_i32 s2, s2, s3
+; GFX7-GISEL-NEXT: v_mov_b32_e32 v0, s2
+; GFX7-GISEL-NEXT: s_mov_b32 s2, -1
+; GFX7-GISEL-NEXT: s_mov_b32 s3, 0xf000
+; GFX7-GISEL-NEXT: buffer_store_byte v0, off, s[0:3], 0
+; GFX7-GISEL-NEXT: s_endpgm
+;
+; GFX12-5-SDAG-LABEL: v_sad_u32_i8_pat1:
+; GFX12-5-SDAG: ; %bb.0:
+; GFX12-5-SDAG-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-SDAG-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-SDAG-NEXT: v_nop
+; GFX12-5-SDAG-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-SDAG-NEXT: s_load_b96 s[0:2], s[4:5], 0x24 nv
+; GFX12-5-SDAG-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-SDAG-NEXT: s_lshr_b32 s3, s2, 8
+; GFX12-5-SDAG-NEXT: s_and_b32 s4, s2, 0xff
+; GFX12-5-SDAG-NEXT: s_and_b32 s3, s3, 0xff
+; GFX12-5-SDAG-NEXT: s_lshr_b32 s2, s2, 16
+; GFX12-5-SDAG-NEXT: s_min_u32 s5, s4, s3
+; GFX12-5-SDAG-NEXT: s_max_u32 s3, s4, s3
+; GFX12-5-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
+; GFX12-5-SDAG-NEXT: s_sub_co_i32 s3, s3, s5
+; GFX12-5-SDAG-NEXT: s_add_co_i32 s2, s3, s2
+; GFX12-5-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
+; GFX12-5-SDAG-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s2
+; GFX12-5-SDAG-NEXT: global_store_b8 v0, v1, s[0:1]
+; GFX12-5-SDAG-NEXT: s_endpgm
+;
+; GFX12-5-GISEL-LABEL: v_sad_u32_i8_pat1:
+; GFX12-5-GISEL: ; %bb.0:
+; GFX12-5-GISEL-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-GISEL-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-GISEL-NEXT: v_nop
+; GFX12-5-GISEL-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-GISEL-NEXT: s_load_b96 s[0:2], s[4:5], 0x24 nv
+; GFX12-5-GISEL-NEXT: v_mov_b32_e32 v1, 0
+; GFX12-5-GISEL-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-GISEL-NEXT: s_lshr_b32 s3, s2, 8
+; GFX12-5-GISEL-NEXT: s_and_b32 s4, s2, 0xff
+; GFX12-5-GISEL-NEXT: s_and_b32 s3, s3, 0xff
+; GFX12-5-GISEL-NEXT: s_lshr_b32 s2, s2, 16
+; GFX12-5-GISEL-NEXT: s_max_u32 s5, s4, s3
+; GFX12-5-GISEL-NEXT: s_min_u32 s3, s4, s3
+; GFX12-5-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
+; GFX12-5-GISEL-NEXT: s_sub_co_i32 s3, s5, s3
+; GFX12-5-GISEL-NEXT: s_add_co_i32 s2, s3, s2
+; GFX12-5-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
+; GFX12-5-GISEL-NEXT: v_mov_b32_e32 v0, s2
+; GFX12-5-GISEL-NEXT: global_store_b8 v1, v0, s[0:1]
+; GFX12-5-GISEL-NEXT: s_endpgm
%icmp0 = icmp ugt i8 %a, %b
%t0 = select i1 %icmp0, i8 %a, i8 %b
@@ -481,24 +1475,105 @@ define amdgpu_kernel void @v_sad_u32_i8_pat1(ptr addrspace(1) %out, i8 %a, i8 %b
}
define amdgpu_kernel void @v_sad_u32_i8_pat2(ptr addrspace(1) %out) {
-; GCN-LABEL: v_sad_u32_i8_pat2:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_add_i32 s12, s12, s17
-; GCN-NEXT: s_mov_b32 flat_scratch_lo, s13
-; GCN-NEXT: s_lshr_b32 flat_scratch_hi, s12, 8
-; GCN-NEXT: flat_load_ubyte v0, v[0:1] glc
-; GCN-NEXT: s_waitcnt vmcnt(0)
-; GCN-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x0
-; GCN-NEXT: flat_load_ubyte v1, v[0:1] glc
-; GCN-NEXT: s_waitcnt vmcnt(0)
-; GCN-NEXT: flat_load_ubyte v2, v[0:1] glc
-; GCN-NEXT: s_waitcnt vmcnt(0)
-; GCN-NEXT: v_sad_u32 v2, v0, v1, v2
-; GCN-NEXT: s_waitcnt lgkmcnt(0)
-; GCN-NEXT: v_mov_b32_e32 v0, s0
-; GCN-NEXT: v_mov_b32_e32 v1, s1
-; GCN-NEXT: flat_store_byte v[0:1], v2
-; GCN-NEXT: s_endpgm
+; GFX7-SDAG-LABEL: v_sad_u32_i8_pat2:
+; GFX7-SDAG: ; %bb.0:
+; GFX7-SDAG-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x9
+; GFX7-SDAG-NEXT: s_mov_b32 s3, 0xf000
+; GFX7-SDAG-NEXT: s_mov_b32 s2, -1
+; GFX7-SDAG-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-SDAG-NEXT: buffer_load_ubyte v0, off, s[0:3], 0 glc
+; GFX7-SDAG-NEXT: s_waitcnt vmcnt(0)
+; GFX7-SDAG-NEXT: buffer_load_ubyte v1, off, s[0:3], 0 glc
+; GFX7-SDAG-NEXT: s_waitcnt vmcnt(0)
+; GFX7-SDAG-NEXT: buffer_load_ubyte v2, off, s[0:3], 0 glc
+; GFX7-SDAG-NEXT: s_waitcnt vmcnt(0)
+; GFX7-SDAG-NEXT: v_sad_u32 v0, v0, v1, v2
+; GFX7-SDAG-NEXT: buffer_store_byte v0, off, s[0:3], 0
+; GFX7-SDAG-NEXT: s_endpgm
+;
+; GFX7-GISEL-LABEL: v_sad_u32_i8_pat2:
+; GFX7-GISEL: ; %bb.0:
+; GFX7-GISEL-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x9
+; GFX7-GISEL-NEXT: s_mov_b32 s2, -1
+; GFX7-GISEL-NEXT: s_mov_b32 s3, 0xf000
+; GFX7-GISEL-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-GISEL-NEXT: buffer_load_ubyte v0, off, s[0:3], 0 glc
+; GFX7-GISEL-NEXT: s_waitcnt vmcnt(0)
+; GFX7-GISEL-NEXT: buffer_load_ubyte v1, off, s[0:3], 0 glc
+; GFX7-GISEL-NEXT: s_waitcnt vmcnt(0)
+; GFX7-GISEL-NEXT: buffer_load_ubyte v2, off, s[0:3], 0 glc
+; GFX7-GISEL-NEXT: s_waitcnt vmcnt(0)
+; GFX7-GISEL-NEXT: v_readfirstlane_b32 s4, v0
+; GFX7-GISEL-NEXT: v_readfirstlane_b32 s5, v1
+; GFX7-GISEL-NEXT: s_cmp_gt_u32 s4, s5
+; GFX7-GISEL-NEXT: s_cselect_b32 s7, 1, 0
+; GFX7-GISEL-NEXT: s_sub_i32 s8, s4, s5
+; GFX7-GISEL-NEXT: s_sub_i32 s4, s5, s4
+; GFX7-GISEL-NEXT: s_cmp_lg_u32 s7, 0
+; GFX7-GISEL-NEXT: v_readfirstlane_b32 s6, v2
+; GFX7-GISEL-NEXT: s_cselect_b32 s4, s8, s4
+; GFX7-GISEL-NEXT: s_add_i32 s4, s4, s6
+; GFX7-GISEL-NEXT: v_mov_b32_e32 v0, s4
+; GFX7-GISEL-NEXT: buffer_store_byte v0, off, s[0:3], 0
+; GFX7-GISEL-NEXT: s_endpgm
+;
+; GFX12-5-SDAG-LABEL: v_sad_u32_i8_pat2:
+; GFX12-5-SDAG: ; %bb.0:
+; GFX12-5-SDAG-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-SDAG-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-SDAG-NEXT: v_nop
+; GFX12-5-SDAG-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-SDAG-NEXT: global_load_u8 v0, v[0:1], off scope:SCOPE_SYS
+; GFX12-5-SDAG-NEXT: s_wait_loadcnt 0x0
+; GFX12-5-SDAG-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv
+; GFX12-5-SDAG-NEXT: s_wait_xcnt 0x0
+; GFX12-5-SDAG-NEXT: global_load_u8 v1, v[0:1], off scope:SCOPE_SYS
+; GFX12-5-SDAG-NEXT: s_wait_loadcnt 0x0
+; GFX12-5-SDAG-NEXT: ; kill: killed $vgpr0_vgpr1
+; GFX12-5-SDAG-NEXT: ; kill: killed $vgpr0_vgpr1
+; GFX12-5-SDAG-NEXT: ; kill: killed $vgpr0_vgpr1
+; GFX12-5-SDAG-NEXT: global_load_u8 v2, v[0:1], off scope:SCOPE_SYS
+; GFX12-5-SDAG-NEXT: s_wait_loadcnt 0x0
+; GFX12-5-SDAG-NEXT: v_min_u32_e32 v3, v0, v1
+; GFX12-5-SDAG-NEXT: v_max_u32_e32 v0, v0, v1
+; GFX12-5-SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
+; GFX12-5-SDAG-NEXT: v_dual_mov_b32 v1, 0 :: v_dual_sub_nc_u32 v0, v0, v3
+; GFX12-5-SDAG-NEXT: v_add_nc_u32_e32 v0, v0, v2
+; GFX12-5-SDAG-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-SDAG-NEXT: global_store_b8 v1, v0, s[0:1]
+; GFX12-5-SDAG-NEXT: s_endpgm
+;
+; GFX12-5-GISEL-LABEL: v_sad_u32_i8_pat2:
+; GFX12-5-GISEL: ; %bb.0:
+; GFX12-5-GISEL-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-GISEL-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-GISEL-NEXT: v_nop
+; GFX12-5-GISEL-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-GISEL-NEXT: global_load_u8 v0, v[0:1], off scope:SCOPE_SYS
+; GFX12-5-GISEL-NEXT: s_wait_loadcnt 0x0
+; GFX12-5-GISEL-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv
+; GFX12-5-GISEL-NEXT: s_wait_xcnt 0x0
+; GFX12-5-GISEL-NEXT: global_load_u8 v1, v[0:1], off scope:SCOPE_SYS
+; GFX12-5-GISEL-NEXT: s_wait_loadcnt 0x0
+; GFX12-5-GISEL-NEXT: v_readfirstlane_b32 s2, v0
+; GFX12-5-GISEL-NEXT: ; kill: killed $vgpr0_vgpr1
+; GFX12-5-GISEL-NEXT: global_load_u8 v2, v[0:1], off scope:SCOPE_SYS
+; GFX12-5-GISEL-NEXT: s_wait_loadcnt 0x0
+; GFX12-5-GISEL-NEXT: v_readfirstlane_b32 s3, v1
+; GFX12-5-GISEL-NEXT: v_mov_b32_e32 v1, 0
+; GFX12-5-GISEL-NEXT: s_cmp_gt_u32 s2, s3
+; GFX12-5-GISEL-NEXT: s_cselect_b32 s4, 1, 0
+; GFX12-5-GISEL-NEXT: s_sub_co_i32 s6, s2, s3
+; GFX12-5-GISEL-NEXT: s_sub_co_i32 s2, s3, s2
+; GFX12-5-GISEL-NEXT: s_cmp_lg_u32 s4, 0
+; GFX12-5-GISEL-NEXT: s_cselect_b32 s2, s6, s2
+; GFX12-5-GISEL-NEXT: v_readfirstlane_b32 s5, v2
+; GFX12-5-GISEL-NEXT: s_add_co_i32 s2, s2, s5
+; GFX12-5-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
+; GFX12-5-GISEL-NEXT: v_mov_b32_e32 v0, s2
+; GFX12-5-GISEL-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-GISEL-NEXT: global_store_b8 v1, v0, s[0:1]
+; GFX12-5-GISEL-NEXT: s_endpgm
%a = load volatile i8, ptr addrspace(1) poison
%b = load volatile i8, ptr addrspace(1) poison
%c = load volatile i8, ptr addrspace(1) poison
@@ -514,24 +1589,90 @@ define amdgpu_kernel void @v_sad_u32_i8_pat2(ptr addrspace(1) %out) {
}
define amdgpu_kernel void @s_sad_u32_i8_pat2(ptr addrspace(1) %out, i8 zeroext %a, i8 zeroext %b, i8 zeroext %c) {
-; GCN-LABEL: s_sad_u32_i8_pat2:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_load_dword s2, s[8:9], 0x2
-; GCN-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x0
-; GCN-NEXT: s_add_i32 s12, s12, s17
-; GCN-NEXT: s_mov_b32 flat_scratch_lo, s13
-; GCN-NEXT: s_lshr_b32 flat_scratch_hi, s12, 8
-; GCN-NEXT: s_waitcnt lgkmcnt(0)
-; GCN-NEXT: s_and_b32 s3, s2, 0xff
-; GCN-NEXT: s_bfe_u32 s4, s2, 0x80008
-; GCN-NEXT: s_lshr_b32 s2, s2, 16
-; GCN-NEXT: v_mov_b32_e32 v0, s4
-; GCN-NEXT: v_mov_b32_e32 v1, s2
-; GCN-NEXT: v_sad_u32 v2, s3, v0, v1
-; GCN-NEXT: v_mov_b32_e32 v0, s0
-; GCN-NEXT: v_mov_b32_e32 v1, s1
-; GCN-NEXT: flat_store_byte v[0:1], v2
-; GCN-NEXT: s_endpgm
+; GFX7-SDAG-LABEL: s_sad_u32_i8_pat2:
+; GFX7-SDAG: ; %bb.0:
+; GFX7-SDAG-NEXT: s_load_dword s6, s[4:5], 0xb
+; GFX7-SDAG-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x9
+; GFX7-SDAG-NEXT: s_mov_b32 s3, 0xf000
+; GFX7-SDAG-NEXT: s_mov_b32 s2, -1
+; GFX7-SDAG-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-SDAG-NEXT: s_and_b32 s4, s6, 0xff
+; GFX7-SDAG-NEXT: s_bfe_u32 s5, s6, 0x80008
+; GFX7-SDAG-NEXT: s_lshr_b32 s6, s6, 16
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v0, s5
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v1, s6
+; GFX7-SDAG-NEXT: v_sad_u32 v0, s4, v0, v1
+; GFX7-SDAG-NEXT: buffer_store_byte v0, off, s[0:3], 0
+; GFX7-SDAG-NEXT: s_endpgm
+;
+; GFX7-GISEL-LABEL: s_sad_u32_i8_pat2:
+; GFX7-GISEL: ; %bb.0:
+; GFX7-GISEL-NEXT: s_load_dword s2, s[4:5], 0xb
+; GFX7-GISEL-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x9
+; GFX7-GISEL-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-GISEL-NEXT: s_lshr_b32 s3, s2, 8
+; GFX7-GISEL-NEXT: s_lshr_b32 s4, s2, 16
+; GFX7-GISEL-NEXT: s_and_b32 s5, s2, 0xff
+; GFX7-GISEL-NEXT: s_and_b32 s6, s3, 0xff
+; GFX7-GISEL-NEXT: s_cmp_gt_u32 s5, s6
+; GFX7-GISEL-NEXT: s_cselect_b32 s5, 1, 0
+; GFX7-GISEL-NEXT: s_sub_i32 s6, s2, s3
+; GFX7-GISEL-NEXT: s_sub_i32 s2, s3, s2
+; GFX7-GISEL-NEXT: s_cmp_lg_u32 s5, 0
+; GFX7-GISEL-NEXT: s_cselect_b32 s2, s6, s2
+; GFX7-GISEL-NEXT: s_add_i32 s2, s2, s4
+; GFX7-GISEL-NEXT: v_mov_b32_e32 v0, s2
+; GFX7-GISEL-NEXT: s_mov_b32 s2, -1
+; GFX7-GISEL-NEXT: s_mov_b32 s3, 0xf000
+; GFX7-GISEL-NEXT: buffer_store_byte v0, off, s[0:3], 0
+; GFX7-GISEL-NEXT: s_endpgm
+;
+; GFX12-5-SDAG-LABEL: s_sad_u32_i8_pat2:
+; GFX12-5-SDAG: ; %bb.0:
+; GFX12-5-SDAG-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-SDAG-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-SDAG-NEXT: v_nop
+; GFX12-5-SDAG-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-SDAG-NEXT: s_load_b96 s[0:2], s[4:5], 0x24 nv
+; GFX12-5-SDAG-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-SDAG-NEXT: s_lshr_b32 s3, s2, 8
+; GFX12-5-SDAG-NEXT: s_and_b32 s4, s2, 0xff
+; GFX12-5-SDAG-NEXT: s_and_b32 s3, s3, 0xff
+; GFX12-5-SDAG-NEXT: s_lshr_b32 s2, s2, 16
+; GFX12-5-SDAG-NEXT: s_min_u32 s5, s4, s3
+; GFX12-5-SDAG-NEXT: s_max_u32 s3, s4, s3
+; GFX12-5-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
+; GFX12-5-SDAG-NEXT: s_sub_co_i32 s3, s3, s5
+; GFX12-5-SDAG-NEXT: s_add_co_i32 s2, s3, s2
+; GFX12-5-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
+; GFX12-5-SDAG-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s2
+; GFX12-5-SDAG-NEXT: global_store_b8 v0, v1, s[0:1]
+; GFX12-5-SDAG-NEXT: s_endpgm
+;
+; GFX12-5-GISEL-LABEL: s_sad_u32_i8_pat2:
+; GFX12-5-GISEL: ; %bb.0:
+; GFX12-5-GISEL-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-GISEL-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-GISEL-NEXT: v_nop
+; GFX12-5-GISEL-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-GISEL-NEXT: s_load_b96 s[0:2], s[4:5], 0x24 nv
+; GFX12-5-GISEL-NEXT: v_mov_b32_e32 v1, 0
+; GFX12-5-GISEL-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-GISEL-NEXT: s_lshr_b32 s3, s2, 8
+; GFX12-5-GISEL-NEXT: s_and_b32 s4, s2, 0xff
+; GFX12-5-GISEL-NEXT: s_and_b32 s5, s3, 0xff
+; GFX12-5-GISEL-NEXT: s_lshr_b32 s6, s2, 16
+; GFX12-5-GISEL-NEXT: s_cmp_gt_u32 s4, s5
+; GFX12-5-GISEL-NEXT: s_cselect_b32 s4, 1, 0
+; GFX12-5-GISEL-NEXT: s_sub_co_i32 s5, s2, s3
+; GFX12-5-GISEL-NEXT: s_sub_co_i32 s2, s3, s2
+; GFX12-5-GISEL-NEXT: s_cmp_lg_u32 s4, 0
+; GFX12-5-GISEL-NEXT: s_cselect_b32 s2, s5, s2
+; GFX12-5-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
+; GFX12-5-GISEL-NEXT: s_add_co_i32 s2, s2, s6
+; GFX12-5-GISEL-NEXT: v_mov_b32_e32 v0, s2
+; GFX12-5-GISEL-NEXT: global_store_b8 v1, v0, s[0:1]
+; GFX12-5-GISEL-NEXT: s_endpgm
%icmp0 = icmp ugt i8 %a, %b
%sub0 = sub i8 %a, %b
%sub1 = sub i8 %b, %a
@@ -544,24 +1685,80 @@ define amdgpu_kernel void @s_sad_u32_i8_pat2(ptr addrspace(1) %out, i8 zeroext %
}
define amdgpu_kernel void @v_sad_u32_mismatched_operands_pat1(ptr addrspace(1) %out, i32 %a, i32 %b, i32 %c, i32 %d) {
-; GCN-LABEL: v_sad_u32_mismatched_operands_pat1:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_load_dwordx4 s[0:3], s[8:9], 0x2
-; GCN-NEXT: s_load_dwordx2 s[4:5], s[8:9], 0x0
-; GCN-NEXT: s_add_i32 s12, s12, s17
-; GCN-NEXT: s_lshr_b32 flat_scratch_hi, s12, 8
-; GCN-NEXT: s_mov_b32 flat_scratch_lo, s13
-; GCN-NEXT: s_waitcnt lgkmcnt(0)
-; GCN-NEXT: s_max_u32 s6, s0, s1
-; GCN-NEXT: s_cmp_le_u32 s0, s1
-; GCN-NEXT: s_cselect_b32 s0, s0, s3
-; GCN-NEXT: s_sub_i32 s0, s6, s0
-; GCN-NEXT: s_add_i32 s0, s0, s2
-; GCN-NEXT: v_mov_b32_e32 v0, s4
-; GCN-NEXT: v_mov_b32_e32 v1, s5
-; GCN-NEXT: v_mov_b32_e32 v2, s0
-; GCN-NEXT: flat_store_dword v[0:1], v2
-; GCN-NEXT: s_endpgm
+; GFX7-SDAG-LABEL: v_sad_u32_mismatched_operands_pat1:
+; GFX7-SDAG: ; %bb.0:
+; GFX7-SDAG-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0xb
+; GFX7-SDAG-NEXT: s_load_dwordx2 s[4:5], s[4:5], 0x9
+; GFX7-SDAG-NEXT: s_mov_b32 s7, 0xf000
+; GFX7-SDAG-NEXT: s_mov_b32 s6, -1
+; GFX7-SDAG-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-SDAG-NEXT: s_max_u32 s8, s0, s1
+; GFX7-SDAG-NEXT: s_cmp_le_u32 s0, s1
+; GFX7-SDAG-NEXT: s_cselect_b32 s0, s0, s3
+; GFX7-SDAG-NEXT: s_sub_i32 s0, s8, s0
+; GFX7-SDAG-NEXT: s_add_i32 s0, s0, s2
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v0, s0
+; GFX7-SDAG-NEXT: buffer_store_dword v0, off, s[4:7], 0
+; GFX7-SDAG-NEXT: s_endpgm
+;
+; GFX7-GISEL-LABEL: v_sad_u32_mismatched_operands_pat1:
+; GFX7-GISEL: ; %bb.0:
+; GFX7-GISEL-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0xb
+; GFX7-GISEL-NEXT: s_load_dwordx2 s[4:5], s[4:5], 0x9
+; GFX7-GISEL-NEXT: s_mov_b32 s6, -1
+; GFX7-GISEL-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-GISEL-NEXT: s_max_u32 s7, s0, s1
+; GFX7-GISEL-NEXT: s_cmp_le_u32 s0, s1
+; GFX7-GISEL-NEXT: s_cselect_b32 s0, s0, s3
+; GFX7-GISEL-NEXT: s_sub_i32 s0, s7, s0
+; GFX7-GISEL-NEXT: s_add_i32 s0, s0, s2
+; GFX7-GISEL-NEXT: v_mov_b32_e32 v0, s0
+; GFX7-GISEL-NEXT: s_mov_b32 s7, 0xf000
+; GFX7-GISEL-NEXT: buffer_store_dword v0, off, s[4:7], 0
+; GFX7-GISEL-NEXT: s_endpgm
+;
+; GFX12-5-SDAG-LABEL: v_sad_u32_mismatched_operands_pat1:
+; GFX12-5-SDAG: ; %bb.0:
+; GFX12-5-SDAG-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-SDAG-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-SDAG-NEXT: v_nop
+; GFX12-5-SDAG-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-SDAG-NEXT: s_clause 0x1
+; GFX12-5-SDAG-NEXT: s_load_b128 s[0:3], s[4:5], 0x2c nv
+; GFX12-5-SDAG-NEXT: s_load_b64 s[6:7], s[4:5], 0x24 nv
+; GFX12-5-SDAG-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-SDAG-NEXT: s_max_u32 s4, s0, s1
+; GFX12-5-SDAG-NEXT: s_cmp_le_u32 s0, s1
+; GFX12-5-SDAG-NEXT: s_cselect_b32 s0, s0, s3
+; GFX12-5-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
+; GFX12-5-SDAG-NEXT: s_sub_co_i32 s0, s4, s0
+; GFX12-5-SDAG-NEXT: s_add_co_i32 s0, s0, s2
+; GFX12-5-SDAG-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
+; GFX12-5-SDAG-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s0
+; GFX12-5-SDAG-NEXT: global_store_b32 v0, v1, s[6:7]
+; GFX12-5-SDAG-NEXT: s_endpgm
+;
+; GFX12-5-GISEL-LABEL: v_sad_u32_mismatched_operands_pat1:
+; GFX12-5-GISEL: ; %bb.0:
+; GFX12-5-GISEL-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-GISEL-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-GISEL-NEXT: v_nop
+; GFX12-5-GISEL-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-GISEL-NEXT: s_clause 0x1
+; GFX12-5-GISEL-NEXT: s_load_b128 s[0:3], s[4:5], 0x2c nv
+; GFX12-5-GISEL-NEXT: s_load_b64 s[6:7], s[4:5], 0x24 nv
+; GFX12-5-GISEL-NEXT: v_mov_b32_e32 v1, 0
+; GFX12-5-GISEL-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-GISEL-NEXT: s_max_u32 s4, s0, s1
+; GFX12-5-GISEL-NEXT: s_cmp_le_u32 s0, s1
+; GFX12-5-GISEL-NEXT: s_cselect_b32 s0, s0, s3
+; GFX12-5-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
+; GFX12-5-GISEL-NEXT: s_sub_co_i32 s0, s4, s0
+; GFX12-5-GISEL-NEXT: s_add_co_i32 s0, s0, s2
+; GFX12-5-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
+; GFX12-5-GISEL-NEXT: v_mov_b32_e32 v0, s0
+; GFX12-5-GISEL-NEXT: global_store_b32 v1, v0, s[6:7]
+; GFX12-5-GISEL-NEXT: s_endpgm
%icmp0 = icmp ugt i32 %a, %b
%t0 = select i1 %icmp0, i32 %a, i32 %b
@@ -576,24 +1773,78 @@ define amdgpu_kernel void @v_sad_u32_mismatched_operands_pat1(ptr addrspace(1) %
}
define amdgpu_kernel void @v_sad_u32_mismatched_operands_pat2(ptr addrspace(1) %out, i32 %a, i32 %b, i32 %c, i32 %d) {
-; GCN-LABEL: v_sad_u32_mismatched_operands_pat2:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_load_dwordx4 s[0:3], s[8:9], 0x2
-; GCN-NEXT: s_load_dwordx2 s[4:5], s[8:9], 0x0
-; GCN-NEXT: s_add_i32 s12, s12, s17
-; GCN-NEXT: s_mov_b32 flat_scratch_lo, s13
-; GCN-NEXT: s_lshr_b32 flat_scratch_hi, s12, 8
-; GCN-NEXT: s_waitcnt lgkmcnt(0)
-; GCN-NEXT: v_mov_b32_e32 v0, s0
-; GCN-NEXT: s_sub_i32 s0, s0, s3
-; GCN-NEXT: v_mov_b32_e32 v1, s0
-; GCN-NEXT: v_sub_i32_e32 v0, vcc, s1, v0
-; GCN-NEXT: v_cndmask_b32_e32 v0, v0, v1, vcc
-; GCN-NEXT: v_add_i32_e32 v2, vcc, s2, v0
-; GCN-NEXT: v_mov_b32_e32 v0, s4
-; GCN-NEXT: v_mov_b32_e32 v1, s5
-; GCN-NEXT: flat_store_dword v[0:1], v2
-; GCN-NEXT: s_endpgm
+; GFX7-SDAG-LABEL: v_sad_u32_mismatched_operands_pat2:
+; GFX7-SDAG: ; %bb.0:
+; GFX7-SDAG-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0xb
+; GFX7-SDAG-NEXT: s_load_dwordx2 s[4:5], s[4:5], 0x9
+; GFX7-SDAG-NEXT: s_mov_b32 s7, 0xf000
+; GFX7-SDAG-NEXT: s_mov_b32 s6, -1
+; GFX7-SDAG-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v0, s0
+; GFX7-SDAG-NEXT: s_sub_i32 s0, s0, s3
+; GFX7-SDAG-NEXT: v_mov_b32_e32 v1, s0
+; GFX7-SDAG-NEXT: v_sub_i32_e32 v0, vcc, s1, v0
+; GFX7-SDAG-NEXT: v_cndmask_b32_e32 v0, v0, v1, vcc
+; GFX7-SDAG-NEXT: v_add_i32_e32 v0, vcc, s2, v0
+; GFX7-SDAG-NEXT: buffer_store_dword v0, off, s[4:7], 0
+; GFX7-SDAG-NEXT: s_endpgm
+;
+; GFX7-GISEL-LABEL: v_sad_u32_mismatched_operands_pat2:
+; GFX7-GISEL: ; %bb.0:
+; GFX7-GISEL-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0xb
+; GFX7-GISEL-NEXT: s_load_dwordx2 s[4:5], s[4:5], 0x9
+; GFX7-GISEL-NEXT: s_mov_b32 s6, -1
+; GFX7-GISEL-NEXT: s_waitcnt lgkmcnt(0)
+; GFX7-GISEL-NEXT: s_sub_u32 s1, s1, s0
+; GFX7-GISEL-NEXT: s_cselect_b32 s7, 1, 0
+; GFX7-GISEL-NEXT: s_sub_i32 s0, s0, s3
+; GFX7-GISEL-NEXT: s_cmp_lg_u32 s7, 0
+; GFX7-GISEL-NEXT: s_cselect_b32 s0, s0, s1
+; GFX7-GISEL-NEXT: s_add_i32 s0, s0, s2
+; GFX7-GISEL-NEXT: v_mov_b32_e32 v0, s0
+; GFX7-GISEL-NEXT: s_mov_b32 s7, 0xf000
+; GFX7-GISEL-NEXT: buffer_store_dword v0, off, s[4:7], 0
+; GFX7-GISEL-NEXT: s_endpgm
+;
+; GFX12-5-SDAG-LABEL: v_sad_u32_mismatched_operands_pat2:
+; GFX12-5-SDAG: ; %bb.0:
+; GFX12-5-SDAG-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-SDAG-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-SDAG-NEXT: v_nop
+; GFX12-5-SDAG-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-SDAG-NEXT: s_clause 0x1
+; GFX12-5-SDAG-NEXT: s_load_b128 s[0:3], s[4:5], 0x2c nv
+; GFX12-5-SDAG-NEXT: s_load_b64 s[6:7], s[4:5], 0x24 nv
+; GFX12-5-SDAG-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-SDAG-NEXT: s_sub_co_i32 s3, s0, s3
+; GFX12-5-SDAG-NEXT: v_sub_co_u32 v0, s0, s1, s0
+; GFX12-5-SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
+; GFX12-5-SDAG-NEXT: v_cndmask_b32_e64 v0, v0, s3, s0
+; GFX12-5-SDAG-NEXT: v_dual_mov_b32 v1, 0 :: v_dual_add_nc_u32 v0, s2, v0
+; GFX12-5-SDAG-NEXT: global_store_b32 v1, v0, s[6:7]
+; GFX12-5-SDAG-NEXT: s_endpgm
+;
+; GFX12-5-GISEL-LABEL: v_sad_u32_mismatched_operands_pat2:
+; GFX12-5-GISEL: ; %bb.0:
+; GFX12-5-GISEL-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GFX12-5-GISEL-NEXT: s_mov_b64 s[64:65], 0
+; GFX12-5-GISEL-NEXT: v_nop
+; GFX12-5-GISEL-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GFX12-5-GISEL-NEXT: s_clause 0x1
+; GFX12-5-GISEL-NEXT: s_load_b128 s[0:3], s[4:5], 0x2c nv
+; GFX12-5-GISEL-NEXT: s_load_b64 s[6:7], s[4:5], 0x24 nv
+; GFX12-5-GISEL-NEXT: v_mov_b32_e32 v1, 0
+; GFX12-5-GISEL-NEXT: s_wait_kmcnt 0x0
+; GFX12-5-GISEL-NEXT: s_sub_co_u32 s1, s1, s0
+; GFX12-5-GISEL-NEXT: s_cselect_b32 s4, 1, 0
+; GFX12-5-GISEL-NEXT: s_sub_co_i32 s0, s0, s3
+; GFX12-5-GISEL-NEXT: s_cmp_lg_u32 s4, 0
+; GFX12-5-GISEL-NEXT: s_cselect_b32 s0, s0, s1
+; GFX12-5-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
+; GFX12-5-GISEL-NEXT: s_add_co_i32 s0, s0, s2
+; GFX12-5-GISEL-NEXT: v_mov_b32_e32 v0, s0
+; GFX12-5-GISEL-NEXT: global_store_b32 v1, v0, s[6:7]
+; GFX12-5-GISEL-NEXT: s_endpgm
%icmp0 = icmp ugt i32 %a, %b
%sub0 = sub i32 %a, %d
%sub1 = sub i32 %b, %a
More information about the llvm-commits
mailing list