[llvm] Fuse mad64_32 from 24-bit multiply-add (PR #225114)
via llvm-commits
llvm-commits at lists.llvm.org
Wed Sep 23 04:02:02 PDT 2026
https://github.com/skc7 updated https://github.com/llvm/llvm-project/pull/225114
>From 17608d625bf29910e9a7cc142cdbf1f791fea1fa Mon Sep 17 00:00:00 2001
From: skc7 <Krishna.Sankisa at amd.com>
Date: Mon, 21 Sep 2026 20:54:08 +0530
Subject: [PATCH] [AMDGPU] Fuse mad64_32 from 24-bit multiply-add
---
llvm/lib/Target/AMDGPU/SIISelLowering.cpp | 106 +++++++
.../atomic_optimizations_global_pointer.ll | 13 +-
.../atomic_optimizations_local_pointer.ll | 13 +-
.../AMDGPU/coexec-mfma-interleave-gfx950.ll | 57 ++--
llvm/test/CodeGen/AMDGPU/idot4-test.ll | 54 ++--
.../CodeGen/AMDGPU/integer-mad-patterns.ll | 279 ++++++++++++------
llvm/test/CodeGen/AMDGPU/mad_u64_u32_mul24.ll | 126 ++++++++
7 files changed, 490 insertions(+), 158 deletions(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/mad_u64_u32_mul24.ll
diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
index f9a4ba4c1bd1d8..f7efd9be523556 100644
--- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
@@ -17440,6 +17440,103 @@ static SDValue tryFoldMADwithSRL(SelectionDAG &DAG, const SDLoc &SL,
DAG.getZeroExtendInReg(AddRHS, SL, MVT::i32), false);
}
+// performMulCombine may lower an i64 ISD::MUL whose operands are known to fit
+// in 24 bits into 24-bit multiply nodes before the enclosing add is combined:
+//
+// mul i64 x, y --> build_pair (mul_u24 x, y), (mulhi_u24 x, y)
+//
+// That rewrite removes the ISD::MUL that tryFoldToMad64_32 keys on, so the
+// fused mad is never formed. Recover it by matching the 24-bit multiply forms
+// directly and folding (add mul24(x, y), z) --> mad_[iu]64_[iu]32 x, y, z.
+static SDValue tryFoldMul24ToMad64_32(SDNode *N, SelectionDAG &DAG,
+ const GCNSubtarget &ST) {
+ assert(N->isAnyAdd());
+
+ EVT VT = N->getValueType(0);
+ if (VT != MVT::i64 || !ST.hasMad64_32())
+ return SDValue();
+
+ // Structurally match a 24-bit multiply that produces an i64 value, returning
+ // its two 32-bit factors and whether it is signed. Handles the build_pair
+ // form getMul24 emits for 64-bit results as well as a bare i64 MUL_[IU]24.
+ auto MatchMul24 = [](SDValue V, SDValue &X, SDValue &Y,
+ bool &Signed) -> bool {
+ if (V.getValueType() != MVT::i64)
+ return false;
+
+ if (V.getOpcode() == ISD::BUILD_PAIR) {
+ SDValue Lo = V.getOperand(0);
+ SDValue Hi = V.getOperand(1);
+ bool IsUnsigned = Lo.getOpcode() == AMDGPUISD::MUL_U24 &&
+ Hi.getOpcode() == AMDGPUISD::MULHI_U24;
+ bool IsSigned = Lo.getOpcode() == AMDGPUISD::MUL_I24 &&
+ Hi.getOpcode() == AMDGPUISD::MULHI_I24;
+ // The halves must be the low/high parts of the same multiply.
+ if ((!IsUnsigned && !IsSigned) || Lo.getOperand(0) != Hi.getOperand(0) ||
+ Lo.getOperand(1) != Hi.getOperand(1))
+ return false;
+ X = Lo.getOperand(0);
+ Y = Lo.getOperand(1);
+ Signed = IsSigned;
+ return true;
+ }
+
+ if (V.getOpcode() == AMDGPUISD::MUL_U24 ||
+ V.getOpcode() == AMDGPUISD::MUL_I24) {
+ X = V.getOperand(0);
+ Y = V.getOperand(1);
+ Signed = V.getOpcode() == AMDGPUISD::MUL_I24;
+ return true;
+ }
+
+ return false;
+ };
+
+ SDValue LHS = N->getOperand(0);
+ SDValue RHS = N->getOperand(1);
+ SDValue X, Y;
+ bool Signed = false;
+ if (!MatchMul24(LHS, X, Y, Signed)) {
+ std::swap(LHS, RHS);
+ if (!MatchMul24(LHS, X, Y, Signed))
+ return SDValue();
+ }
+
+ // Without full-rate 64-bit ops, only fold a single-use multiply to avoid
+ // duplicating it. Checked first so a bail builds no nodes.
+ if (!ST.hasFullRate64Ops()) {
+ bool SingleUse = LHS.hasOneUse();
+ if (SingleUse && LHS.getOpcode() == ISD::BUILD_PAIR)
+ SingleUse =
+ LHS.getOperand(0).hasOneUse() && LHS.getOperand(1).hasOneUse();
+ if (!SingleUse)
+ return SDValue();
+ }
+
+ SDLoc SL(N);
+ X = DAG.getAnyExtOrTrunc(X, SL, MVT::i32);
+ Y = DAG.getAnyExtOrTrunc(Y, SL, MVT::i32);
+ unsigned NumBitsX = Signed ? AMDGPUTargetLowering::numBitsSigned(X, DAG)
+ : AMDGPUTargetLowering::numBitsUnsigned(X, DAG);
+ unsigned NumBitsY = Signed ? AMDGPUTargetLowering::numBitsSigned(Y, DAG)
+ : AMDGPUTargetLowering::numBitsUnsigned(Y, DAG);
+ if (Signed) {
+ // AND cannot preserve sign, so only fold when both operands provably fit.
+ if (NumBitsX > 24 || NumBitsY > 24)
+ return SDValue();
+ } else {
+ // MUL_U24 reads only bits [23:0] but the mad reads all 32. Mask so the
+ // full-width mad computes the same product.
+ SDValue Mask24 = DAG.getConstant((1u << 24) - 1, SL, MVT::i32);
+ if (NumBitsX > 24)
+ X = DAG.getNode(ISD::AND, SL, MVT::i32, X, Mask24);
+ if (NumBitsY > 24)
+ Y = DAG.getNode(ISD::AND, SL, MVT::i32, Y, Mask24);
+ }
+
+ return getMad64_32(DAG, SL, MVT::i64, X, Y, RHS, Signed);
+}
+
// Fold (add (mul x, y), z) --> (mad_[iu]64_[iu]32 x, y, z) plus high
// multiplies, if any.
//
@@ -17880,6 +17977,11 @@ SDValue SITargetLowering::performAddCombine(SDNode *N,
}
}
+ // performMulCombine may have already turned an i64 mul feeding this add into
+ // 24-bit multiply nodes. Recover the fused mad from that form.
+ if (SDValue Folded = tryFoldMul24ToMad64_32(N, DAG, *Subtarget))
+ return Folded;
+
if (SDValue V = reassociateScalarOps(N, DAG)) {
return V;
}
@@ -18111,6 +18213,10 @@ SDValue SITargetLowering::performPtrAddCombine(SDNode *N,
}
}
+ // Analogous recovery for the 24-bit multiply form (see performAddCombine).
+ if (SDValue Folded = tryFoldMul24ToMad64_32(N, DAG, *Subtarget))
+ return Folded;
+
// If the 32 low bits of the constant are all zero, there is nothing to fold
// into an immediate offset, so it's better to eliminate the unnecessary
// addition for the lower 32 bits than to preserve the PTRADD.
diff --git a/llvm/test/CodeGen/AMDGPU/atomic_optimizations_global_pointer.ll b/llvm/test/CodeGen/AMDGPU/atomic_optimizations_global_pointer.ll
index f7045fe08693a1..2b9658cc231ac0 100644
--- a/llvm/test/CodeGen/AMDGPU/atomic_optimizations_global_pointer.ll
+++ b/llvm/test/CodeGen/AMDGPU/atomic_optimizations_global_pointer.ll
@@ -1920,16 +1920,15 @@ define amdgpu_kernel void @add_i64_constant(ptr addrspace(1) %out, ptr addrspace
; GFX8-NEXT: buffer_wbinvl1_vol
; GFX8-NEXT: .LBB3_2:
; GFX8-NEXT: s_or_b64 exec, exec, s[4:5]
-; GFX8-NEXT: v_readfirstlane_b32 s4, v1
-; GFX8-NEXT: v_readfirstlane_b32 s5, v0
-; GFX8-NEXT: v_mul_u32_u24_e32 v0, 5, v2
-; GFX8-NEXT: v_mul_hi_u32_u24_e32 v1, 5, v2
-; GFX8-NEXT: v_mov_b32_e32 v2, s4
-; GFX8-NEXT: v_add_u32_e32 v0, vcc, s5, v0
; GFX8-NEXT: s_waitcnt lgkmcnt(0)
+; GFX8-NEXT: v_readfirstlane_b32 s2, v1
+; GFX8-NEXT: v_readfirstlane_b32 s3, v0
+; GFX8-NEXT: v_mov_b32_e32 v0, s3
+; GFX8-NEXT: v_mov_b32_e32 v1, s2
+; GFX8-NEXT: v_mad_u64_u32 v[0:1], s[2:3], 5, v2, v[0:1]
; GFX8-NEXT: s_mov_b32 s3, 0xf000
; GFX8-NEXT: s_mov_b32 s2, -1
-; GFX8-NEXT: v_addc_u32_e32 v1, vcc, v2, v1, vcc
+; GFX8-NEXT: s_nop 2
; GFX8-NEXT: buffer_store_dwordx2 v[0:1], off, s[0:3], 0
; GFX8-NEXT: s_endpgm
;
diff --git a/llvm/test/CodeGen/AMDGPU/atomic_optimizations_local_pointer.ll b/llvm/test/CodeGen/AMDGPU/atomic_optimizations_local_pointer.ll
index 6f7b956857742c..0ed3f303e2dfb5 100644
--- a/llvm/test/CodeGen/AMDGPU/atomic_optimizations_local_pointer.ll
+++ b/llvm/test/CodeGen/AMDGPU/atomic_optimizations_local_pointer.ll
@@ -1821,17 +1821,16 @@ define amdgpu_kernel void @add_i64_constant(ptr addrspace(1) %out) {
; GFX8-NEXT: s_waitcnt lgkmcnt(0)
; GFX8-NEXT: .LBB4_2:
; GFX8-NEXT: s_or_b64 exec, exec, s[0:1]
+; GFX8-NEXT: v_readfirstlane_b32 s2, v1
+; GFX8-NEXT: v_readfirstlane_b32 s3, v0
; GFX8-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24
-; GFX8-NEXT: v_readfirstlane_b32 s4, v1
-; GFX8-NEXT: v_readfirstlane_b32 s5, v0
-; GFX8-NEXT: v_mul_u32_u24_e32 v0, 5, v2
-; GFX8-NEXT: v_mul_hi_u32_u24_e32 v1, 5, v2
-; GFX8-NEXT: v_mov_b32_e32 v2, s4
-; GFX8-NEXT: v_add_u32_e32 v0, vcc, s5, v0
+; GFX8-NEXT: v_mov_b32_e32 v0, s3
+; GFX8-NEXT: v_mov_b32_e32 v1, s2
+; GFX8-NEXT: v_mad_u64_u32 v[0:1], s[2:3], 5, v2, v[0:1]
; GFX8-NEXT: s_mov_b32 s3, 0xf000
; GFX8-NEXT: s_mov_b32 s2, -1
-; GFX8-NEXT: v_addc_u32_e32 v1, vcc, v2, v1, vcc
; GFX8-NEXT: s_waitcnt lgkmcnt(0)
+; GFX8-NEXT: s_nop 1
; GFX8-NEXT: buffer_store_dwordx2 v[0:1], off, s[0:3], 0
; GFX8-NEXT: s_endpgm
;
diff --git a/llvm/test/CodeGen/AMDGPU/coexec-mfma-interleave-gfx950.ll b/llvm/test/CodeGen/AMDGPU/coexec-mfma-interleave-gfx950.ll
index 3ebfd4d1d5a24c..24b20575b4dfa2 100644
--- a/llvm/test/CodeGen/AMDGPU/coexec-mfma-interleave-gfx950.ll
+++ b/llvm/test/CodeGen/AMDGPU/coexec-mfma-interleave-gfx950.ll
@@ -11,44 +11,43 @@ define amdgpu_kernel void @mfma_16x16_interleave(
; CHECK-LABEL: mfma_16x16_interleave:
; CHECK: ; %bb.0:
; CHECK-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x0
+; CHECK-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x8
; CHECK-NEXT: v_and_b32_e32 v14, 0x3ff, v0
-; CHECK-NEXT: v_mov_b32_e32 v1, 0
-; CHECK-NEXT: v_lshlrev_b32_e32 v0, 4, v14
-; CHECK-NEXT: v_mul_hi_i32_i24_e32 v15, -12, v14
-; CHECK-NEXT: v_mul_i32_i24_e32 v14, -12, v14
+; CHECK-NEXT: v_mov_b32_e32 v29, 0
+; CHECK-NEXT: v_lshlrev_b32_e32 v28, 4, v14
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
-; CHECK-NEXT: v_lshl_add_u64 v[16:17], s[0:1], 0, v[0:1]
-; CHECK-NEXT: global_load_dwordx4 v[2:5], v0, s[0:1]
-; CHECK-NEXT: global_load_dwordx4 v[6:9], v0, s[0:1] offset:16
-; CHECK-NEXT: global_load_dwordx4 v[10:13], v0, s[0:1] offset:32
-; CHECK-NEXT: v_lshl_add_u64 v[18:19], v[16:17], 0, v[14:15]
-; CHECK-NEXT: global_load_dwordx2 v[30:31], v[18:19], off offset:16
-; CHECK-NEXT: global_load_dwordx4 v[14:17], v[18:19], off
-; CHECK-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x8
+; CHECK-NEXT: v_lshl_add_u64 v[12:13], s[0:1], 0, v[28:29]
+; CHECK-NEXT: global_load_dwordx4 v[0:3], v28, s[0:1]
+; CHECK-NEXT: global_load_dwordx4 v[4:7], v28, s[0:1] offset:16
+; CHECK-NEXT: global_load_dwordx4 v[8:11], v28, s[0:1] offset:32
+; CHECK-NEXT: v_mad_i64_i32 v[16:17], s[0:1], v14, -12, v[12:13]
+; CHECK-NEXT: global_load_dwordx2 v[30:31], v[16:17], off offset:16
+; CHECK-NEXT: global_load_dwordx4 v[12:15], v[16:17], off
; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; CHECK-NEXT: ; sched_barrier mask(0x00000000)
-; CHECK-NEXT: v_mfma_f32_16x16x32_f16 v[18:21], v[2:5], v[6:9], v[2:5]
+; CHECK-NEXT: v_mfma_f32_16x16x32_f16 v[16:19], v[0:3], v[4:7], v[0:3]
; CHECK-NEXT: s_add_i32 s0, s8, s9
-; CHECK-NEXT: v_add_u32_e32 v22, v14, v15
-; CHECK-NEXT: v_add_u32_e32 v23, v15, v16
-; CHECK-NEXT: v_mfma_f32_16x16x32_f16 v[26:29], v[2:5], v[6:9], v[6:9]
+; CHECK-NEXT: v_add_u32_e32 v20, v12, v13
+; CHECK-NEXT: v_add_u32_e32 v21, v13, v14
+; CHECK-NEXT: v_mfma_f32_16x16x32_f16 v[24:27], v[0:3], v[4:7], v[4:7]
; CHECK-NEXT: s_add_i32 s1, s10, s11
-; CHECK-NEXT: v_add_u32_e32 v24, v16, v17
-; CHECK-NEXT: v_add_u32_e32 v25, v17, v30
-; CHECK-NEXT: v_mfma_f32_16x16x32_f16 v[2:5], v[2:5], v[6:9], v[10:13]
+; CHECK-NEXT: v_add_u32_e32 v22, v14, v15
+; CHECK-NEXT: v_add_u32_e32 v23, v15, v30
+; CHECK-NEXT: v_mfma_f32_16x16x32_f16 v[0:3], v[0:3], v[4:7], v[8:11]
; CHECK-NEXT: s_add_i32 s2, s12, s13
-; CHECK-NEXT: v_add_u32_e32 v6, v30, v31
-; CHECK-NEXT: v_add_u32_e32 v7, v31, v14
+; CHECK-NEXT: v_add_u32_e32 v4, v30, v31
+; CHECK-NEXT: v_add_u32_e32 v5, v31, v12
; CHECK-NEXT: ; sched_barrier mask(0x00000000)
-; CHECK-NEXT: v_mov_b32_e32 v8, s0
-; CHECK-NEXT: v_mov_b32_e32 v9, s1
+; CHECK-NEXT: global_store_dwordx4 v29, v[16:19], s[14:15]
+; CHECK-NEXT: global_store_dwordx4 v29, v[24:27], s[14:15] offset:16
+; CHECK-NEXT: s_nop 2
+; CHECK-NEXT: global_store_dwordx4 v29, v[0:3], s[14:15] offset:32
+; CHECK-NEXT: global_store_dwordx4 v29, v[20:23], s[14:15] offset:48
+; CHECK-NEXT: v_mov_b32_e32 v6, s0
+; CHECK-NEXT: v_mov_b32_e32 v7, s1
; CHECK-NEXT: v_mov_b32_e32 v0, s2
-; CHECK-NEXT: global_store_dwordx4 v1, v[18:21], s[14:15]
-; CHECK-NEXT: global_store_dwordx4 v1, v[26:29], s[14:15] offset:16
-; CHECK-NEXT: global_store_dwordx4 v1, v[2:5], s[14:15] offset:32
-; CHECK-NEXT: global_store_dwordx4 v1, v[22:25], s[14:15] offset:48
-; CHECK-NEXT: global_store_dwordx4 v1, v[6:9], s[14:15] offset:64
-; CHECK-NEXT: global_store_dword v1, v0, s[14:15] offset:80
+; CHECK-NEXT: global_store_dwordx4 v29, v[4:7], s[14:15] offset:64
+; CHECK-NEXT: global_store_dword v29, v0, s[14:15] offset:80
; CHECK-NEXT: s_endpgm
ptr addrspace(1) %ptr,
i32 %s0, i32 %s1, i32 %s2, i32 %s3,
diff --git a/llvm/test/CodeGen/AMDGPU/idot4-test.ll b/llvm/test/CodeGen/AMDGPU/idot4-test.ll
index 79017f3ea3cd77..9b43197e2e8c85 100644
--- a/llvm/test/CodeGen/AMDGPU/idot4-test.ll
+++ b/llvm/test/CodeGen/AMDGPU/idot4-test.ll
@@ -1175,51 +1175,45 @@ define i64 @dot4_acc64_not_dot4(i8 %a0, i8 %a1, i8 %b0, i8 %b1, i64 %acc) {
; GFX9-DL-LABEL: dot4_acc64_not_dot4:
; GFX9-DL: ; %bb.0:
; GFX9-DL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GFX9-DL-NEXT: v_mul_hi_i32_i24_sdwa v6, sext(v0), sext(v2) dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_0 src1_sel:BYTE_0
-; GFX9-DL-NEXT: v_mul_i32_i24_sdwa v0, sext(v0), sext(v2) dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_0 src1_sel:BYTE_0
-; GFX9-DL-NEXT: v_mul_hi_i32_i24_sdwa v2, sext(v1), sext(v3) dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_0 src1_sel:BYTE_0
-; GFX9-DL-NEXT: v_mul_i32_i24_sdwa v1, sext(v1), sext(v3) dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_0 src1_sel:BYTE_0
-; GFX9-DL-NEXT: v_add_co_u32_e32 v1, vcc, v1, v4
-; GFX9-DL-NEXT: v_addc_co_u32_e32 v2, vcc, v2, v5, vcc
-; GFX9-DL-NEXT: v_add_co_u32_e32 v0, vcc, v0, v1
-; GFX9-DL-NEXT: v_addc_co_u32_e32 v1, vcc, v6, v2, vcc
+; GFX9-DL-NEXT: v_bfe_i32 v1, v1, 0, 8
+; GFX9-DL-NEXT: v_bfe_i32 v3, v3, 0, 8
+; GFX9-DL-NEXT: v_mad_i64_i32 v[3:4], s[4:5], v1, v3, v[4:5]
+; GFX9-DL-NEXT: v_bfe_i32 v0, v0, 0, 8
+; GFX9-DL-NEXT: v_bfe_i32 v1, v2, 0, 8
+; GFX9-DL-NEXT: v_mad_i64_i32 v[0:1], s[4:5], v0, v1, v[3:4]
; GFX9-DL-NEXT: s_setpc_b64 s[30:31]
;
; GFX10-DL-LABEL: dot4_acc64_not_dot4:
; GFX10-DL: ; %bb.0:
; GFX10-DL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GFX10-DL-NEXT: v_mul_i32_i24_sdwa v6, sext(v1), sext(v3) dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_0 src1_sel:BYTE_0
-; GFX10-DL-NEXT: v_mul_hi_i32_i24_sdwa v1, sext(v1), sext(v3) dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_0 src1_sel:BYTE_0
-; GFX10-DL-NEXT: v_mul_hi_i32_i24_sdwa v7, sext(v0), sext(v2) dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_0 src1_sel:BYTE_0
-; GFX10-DL-NEXT: v_mul_i32_i24_sdwa v0, sext(v0), sext(v2) dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_0 src1_sel:BYTE_0
-; GFX10-DL-NEXT: v_add_co_u32 v2, vcc_lo, v6, v4
-; GFX10-DL-NEXT: v_add_co_ci_u32_e32 v1, vcc_lo, v1, v5, vcc_lo
-; GFX10-DL-NEXT: v_add_co_u32 v0, vcc_lo, v0, v2
-; GFX10-DL-NEXT: v_add_co_ci_u32_e32 v1, vcc_lo, v7, v1, vcc_lo
+; GFX10-DL-NEXT: v_bfe_i32 v1, v1, 0, 8
+; GFX10-DL-NEXT: v_bfe_i32 v3, v3, 0, 8
+; GFX10-DL-NEXT: v_bfe_i32 v0, v0, 0, 8
+; GFX10-DL-NEXT: v_mad_i64_i32 v[3:4], s4, v1, v3, v[4:5]
+; GFX10-DL-NEXT: v_bfe_i32 v1, v2, 0, 8
+; GFX10-DL-NEXT: v_mad_i64_i32 v[0:1], s4, v0, v1, v[3:4]
; GFX10-DL-NEXT: s_setpc_b64 s[30:31]
;
; GFX950-LABEL: dot4_acc64_not_dot4:
; GFX950: ; %bb.0:
; GFX950-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GFX950-NEXT: v_mul_hi_i32_i24_sdwa v9, sext(v1), sext(v3) dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_0 src1_sel:BYTE_0
-; GFX950-NEXT: v_mul_i32_i24_sdwa v8, sext(v1), sext(v3) dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_0 src1_sel:BYTE_0
-; GFX950-NEXT: v_mul_hi_i32_i24_sdwa v7, sext(v0), sext(v2) dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_0 src1_sel:BYTE_0
-; GFX950-NEXT: v_mul_i32_i24_sdwa v6, sext(v0), sext(v2) dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_0 src1_sel:BYTE_0
-; GFX950-NEXT: v_lshl_add_u64 v[0:1], v[8:9], 0, v[4:5]
-; GFX950-NEXT: v_lshl_add_u64 v[0:1], v[6:7], 0, v[0:1]
+; GFX950-NEXT: v_bfe_i32 v6, v0, 0, 8
+; GFX950-NEXT: v_bfe_i32 v0, v1, 0, 8
+; GFX950-NEXT: v_bfe_i32 v1, v3, 0, 8
+; GFX950-NEXT: v_bfe_i32 v2, v2, 0, 8
+; GFX950-NEXT: v_mad_i64_i32 v[0:1], s[0:1], v0, v1, v[4:5]
+; GFX950-NEXT: v_mad_i64_i32 v[0:1], s[0:1], v6, v2, v[0:1]
; GFX950-NEXT: s_setpc_b64 s[30:31]
;
; GFX9-NODL-LABEL: dot4_acc64_not_dot4:
; GFX9-NODL: ; %bb.0:
; GFX9-NODL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GFX9-NODL-NEXT: v_mul_hi_i32_i24_sdwa v6, sext(v0), sext(v2) dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_0 src1_sel:BYTE_0
-; GFX9-NODL-NEXT: v_mul_i32_i24_sdwa v0, sext(v0), sext(v2) dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_0 src1_sel:BYTE_0
-; GFX9-NODL-NEXT: v_mul_hi_i32_i24_sdwa v2, sext(v1), sext(v3) dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_0 src1_sel:BYTE_0
-; GFX9-NODL-NEXT: v_mul_i32_i24_sdwa v1, sext(v1), sext(v3) dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_0 src1_sel:BYTE_0
-; GFX9-NODL-NEXT: v_add_co_u32_e32 v1, vcc, v1, v4
-; GFX9-NODL-NEXT: v_addc_co_u32_e32 v2, vcc, v2, v5, vcc
-; GFX9-NODL-NEXT: v_add_co_u32_e32 v0, vcc, v0, v1
-; GFX9-NODL-NEXT: v_addc_co_u32_e32 v1, vcc, v6, v2, vcc
+; GFX9-NODL-NEXT: v_bfe_i32 v1, v1, 0, 8
+; GFX9-NODL-NEXT: v_bfe_i32 v3, v3, 0, 8
+; GFX9-NODL-NEXT: v_mad_i64_i32 v[3:4], s[4:5], v1, v3, v[4:5]
+; GFX9-NODL-NEXT: v_bfe_i32 v0, v0, 0, 8
+; GFX9-NODL-NEXT: v_bfe_i32 v1, v2, 0, 8
+; GFX9-NODL-NEXT: v_mad_i64_i32 v[0:1], s[4:5], v0, v1, v[3:4]
; GFX9-NODL-NEXT: s_setpc_b64 s[30:31]
%cv1e0 = sext i8 %a0 to i64
%cv1e1 = sext i8 %a1 to i64
diff --git a/llvm/test/CodeGen/AMDGPU/integer-mad-patterns.ll b/llvm/test/CodeGen/AMDGPU/integer-mad-patterns.ll
index 3b0f6308771fb5..0c951fe87f9513 100644
--- a/llvm/test/CodeGen/AMDGPU/integer-mad-patterns.ll
+++ b/llvm/test/CodeGen/AMDGPU/integer-mad-patterns.ll
@@ -12466,76 +12466,159 @@ entry:
}
define i64 @mul_u24_add64(i32 %x, i32 %y, i64 %z) {
-; GFX67-LABEL: mul_u24_add64:
-; GFX67: ; %bb.0:
-; GFX67-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GFX67-NEXT: v_mul_hi_u32_u24_e32 v4, v0, v1
-; GFX67-NEXT: v_mul_u32_u24_e32 v0, v0, v1
-; GFX67-NEXT: v_add_i32_e32 v0, vcc, v0, v2
-; GFX67-NEXT: v_addc_u32_e32 v1, vcc, v4, v3, vcc
-; GFX67-NEXT: s_setpc_b64 s[30:31]
+; GFX6-LABEL: mul_u24_add64:
+; GFX6: ; %bb.0:
+; GFX6-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX6-NEXT: v_mul_hi_u32_u24_e32 v4, v0, v1
+; GFX6-NEXT: v_mul_u32_u24_e32 v0, v0, v1
+; GFX6-NEXT: v_add_i32_e32 v0, vcc, v0, v2
+; GFX6-NEXT: v_addc_u32_e32 v1, vcc, v4, v3, vcc
+; GFX6-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX7-SDAG-LABEL: mul_u24_add64:
+; GFX7-SDAG: ; %bb.0:
+; GFX7-SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX7-SDAG-NEXT: v_and_b32_e32 v1, 0xffffff, v1
+; GFX7-SDAG-NEXT: v_and_b32_e32 v0, 0xffffff, v0
+; GFX7-SDAG-NEXT: v_mad_u64_u32 v[0:1], s[4:5], v0, v1, v[2:3]
+; GFX7-SDAG-NEXT: s_setpc_b64 s[30:31]
;
-; GFX8-LABEL: mul_u24_add64:
-; GFX8: ; %bb.0:
-; GFX8-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GFX8-NEXT: v_mul_hi_u32_u24_e32 v4, v0, v1
-; GFX8-NEXT: v_mul_u32_u24_e32 v0, v0, v1
-; GFX8-NEXT: v_add_u32_e32 v0, vcc, v0, v2
-; GFX8-NEXT: v_addc_u32_e32 v1, vcc, v4, v3, vcc
-; GFX8-NEXT: s_setpc_b64 s[30:31]
+; GFX7-GISEL-LABEL: mul_u24_add64:
+; GFX7-GISEL: ; %bb.0:
+; GFX7-GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX7-GISEL-NEXT: v_mul_hi_u32_u24_e32 v4, v0, v1
+; GFX7-GISEL-NEXT: v_mul_u32_u24_e32 v0, v0, v1
+; GFX7-GISEL-NEXT: v_add_i32_e32 v0, vcc, v0, v2
+; GFX7-GISEL-NEXT: v_addc_u32_e32 v1, vcc, v4, v3, vcc
+; GFX7-GISEL-NEXT: s_setpc_b64 s[30:31]
;
-; GFX9-LABEL: mul_u24_add64:
-; GFX9: ; %bb.0:
-; GFX9-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GFX9-NEXT: v_mul_hi_u32_u24_e32 v4, v0, v1
-; GFX9-NEXT: v_mul_u32_u24_e32 v0, v0, v1
-; GFX9-NEXT: v_add_co_u32_e32 v0, vcc, v0, v2
-; GFX9-NEXT: v_addc_co_u32_e32 v1, vcc, v4, v3, vcc
-; GFX9-NEXT: s_setpc_b64 s[30:31]
+; GFX8-SDAG-LABEL: mul_u24_add64:
+; GFX8-SDAG: ; %bb.0:
+; GFX8-SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX8-SDAG-NEXT: v_and_b32_e32 v1, 0xffffff, v1
+; GFX8-SDAG-NEXT: v_and_b32_e32 v0, 0xffffff, v0
+; GFX8-SDAG-NEXT: v_mad_u64_u32 v[0:1], s[4:5], v0, v1, v[2:3]
+; GFX8-SDAG-NEXT: s_setpc_b64 s[30:31]
;
-; GFX10-LABEL: mul_u24_add64:
-; GFX10: ; %bb.0:
-; GFX10-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GFX10-NEXT: v_mul_u32_u24_e32 v4, v0, v1
-; GFX10-NEXT: v_mul_hi_u32_u24_e32 v1, v0, v1
-; GFX10-NEXT: v_add_co_u32 v0, vcc_lo, v4, v2
-; GFX10-NEXT: v_add_co_ci_u32_e64 v1, null, v1, v3, vcc_lo
-; GFX10-NEXT: s_setpc_b64 s[30:31]
+; GFX8-GISEL-LABEL: mul_u24_add64:
+; GFX8-GISEL: ; %bb.0:
+; GFX8-GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX8-GISEL-NEXT: v_mul_hi_u32_u24_e32 v4, v0, v1
+; GFX8-GISEL-NEXT: v_mul_u32_u24_e32 v0, v0, v1
+; GFX8-GISEL-NEXT: v_add_u32_e32 v0, vcc, v0, v2
+; GFX8-GISEL-NEXT: v_addc_u32_e32 v1, vcc, v4, v3, vcc
+; GFX8-GISEL-NEXT: s_setpc_b64 s[30:31]
;
-; GFX11-LABEL: mul_u24_add64:
-; GFX11: ; %bb.0:
-; GFX11-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GFX11-NEXT: v_mul_u32_u24_e32 v4, v0, v1
-; GFX11-NEXT: v_mul_hi_u32_u24_e32 v1, v0, v1
-; GFX11-NEXT: s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
-; GFX11-NEXT: v_add_co_u32 v0, vcc_lo, v4, v2
-; GFX11-NEXT: v_add_co_ci_u32_e64 v1, null, v1, v3, vcc_lo
-; GFX11-NEXT: s_setpc_b64 s[30:31]
+; GFX9-SDAG-LABEL: mul_u24_add64:
+; GFX9-SDAG: ; %bb.0:
+; GFX9-SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX9-SDAG-NEXT: v_and_b32_e32 v1, 0xffffff, v1
+; GFX9-SDAG-NEXT: v_and_b32_e32 v0, 0xffffff, v0
+; GFX9-SDAG-NEXT: v_mad_u64_u32 v[0:1], s[4:5], v0, v1, v[2:3]
+; GFX9-SDAG-NEXT: s_setpc_b64 s[30:31]
;
-; GFX1200-LABEL: mul_u24_add64:
-; GFX1200: ; %bb.0:
-; GFX1200-NEXT: s_wait_loadcnt_dscnt 0x0
-; GFX1200-NEXT: s_wait_expcnt 0x0
-; GFX1200-NEXT: s_wait_samplecnt 0x0
-; GFX1200-NEXT: s_wait_bvhcnt 0x0
-; GFX1200-NEXT: s_wait_kmcnt 0x0
-; GFX1200-NEXT: v_mul_u32_u24_e32 v4, v0, v1
-; GFX1200-NEXT: v_mul_hi_u32_u24_e32 v1, v0, v1
-; GFX1200-NEXT: s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
-; GFX1200-NEXT: v_add_co_u32 v0, vcc_lo, v4, v2
-; GFX1200-NEXT: s_wait_alu depctr_va_vcc(0)
-; GFX1200-NEXT: v_add_co_ci_u32_e64 v1, null, v1, v3, vcc_lo
-; GFX1200-NEXT: s_setpc_b64 s[30:31]
+; GFX9-GISEL-LABEL: mul_u24_add64:
+; GFX9-GISEL: ; %bb.0:
+; GFX9-GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX9-GISEL-NEXT: v_mul_hi_u32_u24_e32 v4, v0, v1
+; GFX9-GISEL-NEXT: v_mul_u32_u24_e32 v0, v0, v1
+; GFX9-GISEL-NEXT: v_add_co_u32_e32 v0, vcc, v0, v2
+; GFX9-GISEL-NEXT: v_addc_co_u32_e32 v1, vcc, v4, v3, vcc
+; GFX9-GISEL-NEXT: s_setpc_b64 s[30:31]
;
-; GFX1250-LABEL: mul_u24_add64:
-; GFX1250: ; %bb.0:
-; GFX1250-NEXT: s_wait_loadcnt_dscnt 0x0
-; GFX1250-NEXT: s_wait_kmcnt 0x0
-; GFX1250-NEXT: v_mul_hi_u32_u24_e32 v5, v0, v1
-; GFX1250-NEXT: v_mul_u32_u24_e32 v4, v0, v1
-; GFX1250-NEXT: s_delay_alu instid0(VALU_DEP_1)
-; GFX1250-NEXT: v_add_nc_u64_e32 v[0:1], v[4:5], v[2:3]
-; GFX1250-NEXT: s_set_pc_i64 s[30:31]
+; GFX10-SDAG-LABEL: mul_u24_add64:
+; GFX10-SDAG: ; %bb.0:
+; GFX10-SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX10-SDAG-NEXT: v_and_b32_e32 v1, 0xffffff, v1
+; GFX10-SDAG-NEXT: v_and_b32_e32 v0, 0xffffff, v0
+; GFX10-SDAG-NEXT: v_mad_u64_u32 v[0:1], null, v0, v1, v[2:3]
+; GFX10-SDAG-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX10-GISEL-LABEL: mul_u24_add64:
+; GFX10-GISEL: ; %bb.0:
+; GFX10-GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX10-GISEL-NEXT: v_mul_u32_u24_e32 v4, v0, v1
+; GFX10-GISEL-NEXT: v_mul_hi_u32_u24_e32 v1, v0, v1
+; GFX10-GISEL-NEXT: v_add_co_u32 v0, vcc_lo, v4, v2
+; GFX10-GISEL-NEXT: v_add_co_ci_u32_e64 v1, null, v1, v3, vcc_lo
+; GFX10-GISEL-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX11-SDAG-LABEL: mul_u24_add64:
+; GFX11-SDAG: ; %bb.0:
+; GFX11-SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX11-SDAG-NEXT: v_and_b32_e32 v4, 0xffffff, v1
+; GFX11-SDAG-NEXT: v_and_b32_e32 v5, 0xffffff, v0
+; GFX11-SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX11-SDAG-NEXT: v_mad_u64_u32 v[0:1], null, v5, v4, v[2:3]
+; GFX11-SDAG-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX11-GISEL-LABEL: mul_u24_add64:
+; GFX11-GISEL: ; %bb.0:
+; GFX11-GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX11-GISEL-NEXT: v_mul_u32_u24_e32 v4, v0, v1
+; GFX11-GISEL-NEXT: v_mul_hi_u32_u24_e32 v1, v0, v1
+; GFX11-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
+; GFX11-GISEL-NEXT: v_add_co_u32 v0, vcc_lo, v4, v2
+; GFX11-GISEL-NEXT: v_add_co_ci_u32_e64 v1, null, v1, v3, vcc_lo
+; GFX11-GISEL-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX1200-SDAG-LABEL: mul_u24_add64:
+; GFX1200-SDAG: ; %bb.0:
+; GFX1200-SDAG-NEXT: s_wait_loadcnt_dscnt 0x0
+; GFX1200-SDAG-NEXT: s_wait_expcnt 0x0
+; GFX1200-SDAG-NEXT: s_wait_samplecnt 0x0
+; GFX1200-SDAG-NEXT: s_wait_bvhcnt 0x0
+; GFX1200-SDAG-NEXT: s_wait_kmcnt 0x0
+; GFX1200-SDAG-NEXT: v_and_b32_e32 v1, 0xffffff, v1
+; GFX1200-SDAG-NEXT: v_and_b32_e32 v0, 0xffffff, v0
+; GFX1200-SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX1200-SDAG-NEXT: v_mad_co_u64_u32 v[0:1], null, v0, v1, v[2:3]
+; GFX1200-SDAG-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX1200-GISEL-LABEL: mul_u24_add64:
+; GFX1200-GISEL: ; %bb.0:
+; GFX1200-GISEL-NEXT: s_wait_loadcnt_dscnt 0x0
+; GFX1200-GISEL-NEXT: s_wait_expcnt 0x0
+; GFX1200-GISEL-NEXT: s_wait_samplecnt 0x0
+; GFX1200-GISEL-NEXT: s_wait_bvhcnt 0x0
+; GFX1200-GISEL-NEXT: s_wait_kmcnt 0x0
+; GFX1200-GISEL-NEXT: v_mul_u32_u24_e32 v4, v0, v1
+; GFX1200-GISEL-NEXT: v_mul_hi_u32_u24_e32 v1, v0, v1
+; GFX1200-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
+; GFX1200-GISEL-NEXT: v_add_co_u32 v0, vcc_lo, v4, v2
+; GFX1200-GISEL-NEXT: s_wait_alu depctr_va_vcc(0)
+; GFX1200-GISEL-NEXT: v_add_co_ci_u32_e64 v1, null, v1, v3, vcc_lo
+; GFX1200-GISEL-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX1250-SDAG-LABEL: mul_u24_add64:
+; GFX1250-SDAG: ; %bb.0:
+; GFX1250-SDAG-NEXT: s_wait_loadcnt_dscnt 0x0
+; GFX1250-SDAG-NEXT: s_wait_kmcnt 0x0
+; GFX1250-SDAG-NEXT: v_and_b32_e32 v1, 0xffffff, v1
+; GFX1250-SDAG-NEXT: v_and_b32_e32 v0, 0xffffff, v0
+; GFX1250-SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX1250-SDAG-NEXT: v_mad_nc_u64_u32 v[0:1], v0, v1, v[2:3]
+; GFX1250-SDAG-NEXT: s_set_pc_i64 s[30:31]
+;
+; GFX1250-GISEL-FAKE16-LABEL: mul_u24_add64:
+; GFX1250-GISEL-FAKE16: ; %bb.0:
+; GFX1250-GISEL-FAKE16-NEXT: s_wait_loadcnt_dscnt 0x0
+; GFX1250-GISEL-FAKE16-NEXT: s_wait_kmcnt 0x0
+; GFX1250-GISEL-FAKE16-NEXT: v_mul_hi_u32_u24_e32 v5, v0, v1
+; GFX1250-GISEL-FAKE16-NEXT: v_mul_u32_u24_e32 v4, v0, v1
+; GFX1250-GISEL-FAKE16-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX1250-GISEL-FAKE16-NEXT: v_add_nc_u64_e32 v[0:1], v[4:5], v[2:3]
+; GFX1250-GISEL-FAKE16-NEXT: s_set_pc_i64 s[30:31]
+;
+; GFX1250-GISEL-REAL16-LABEL: mul_u24_add64:
+; GFX1250-GISEL-REAL16: ; %bb.0:
+; GFX1250-GISEL-REAL16-NEXT: s_wait_loadcnt_dscnt 0x0
+; GFX1250-GISEL-REAL16-NEXT: s_wait_kmcnt 0x0
+; GFX1250-GISEL-REAL16-NEXT: v_mul_hi_u32_u24_e32 v5, v0, v1
+; GFX1250-GISEL-REAL16-NEXT: v_mul_u32_u24_e32 v4, v0, v1
+; GFX1250-GISEL-REAL16-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX1250-GISEL-REAL16-NEXT: v_add_nc_u64_e32 v[0:1], v[4:5], v[2:3]
+; GFX1250-GISEL-REAL16-NEXT: s_set_pc_i64 s[30:31]
%mul = call i64 @llvm.amdgcn.mul.u24.i64(i32 %x, i32 %y)
%add = add i64 %mul, %z
ret i64 %add
@@ -12613,25 +12696,52 @@ define i64 @mul_u24_zext_add64(i32 %x, i32 %y, i64 %z) {
}
define i64 @mul_u24_known_24bit_add64(i16 %x.s, i16 %y.s, i64 %z) {
-; GFX67-LABEL: mul_u24_known_24bit_add64:
-; GFX67: ; %bb.0:
-; GFX67-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GFX67-NEXT: v_and_b32_e32 v0, 0xffff, v0
-; GFX67-NEXT: v_and_b32_e32 v1, 0xffff, v1
-; GFX67-NEXT: v_mul_hi_u32_u24_e32 v4, v0, v1
-; GFX67-NEXT: v_mul_u32_u24_e32 v0, v0, v1
-; GFX67-NEXT: v_add_i32_e32 v0, vcc, v0, v2
-; GFX67-NEXT: v_addc_u32_e32 v1, vcc, v4, v3, vcc
-; GFX67-NEXT: s_setpc_b64 s[30:31]
+; GFX6-LABEL: mul_u24_known_24bit_add64:
+; GFX6: ; %bb.0:
+; GFX6-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX6-NEXT: v_and_b32_e32 v0, 0xffff, v0
+; GFX6-NEXT: v_and_b32_e32 v1, 0xffff, v1
+; GFX6-NEXT: v_mul_hi_u32_u24_e32 v4, v0, v1
+; GFX6-NEXT: v_mul_u32_u24_e32 v0, v0, v1
+; GFX6-NEXT: v_add_i32_e32 v0, vcc, v0, v2
+; GFX6-NEXT: v_addc_u32_e32 v1, vcc, v4, v3, vcc
+; GFX6-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX7-SDAG-LABEL: mul_u24_known_24bit_add64:
+; GFX7-SDAG: ; %bb.0:
+; GFX7-SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX7-SDAG-NEXT: v_and_b32_e32 v0, 0xffff, v0
+; GFX7-SDAG-NEXT: v_and_b32_e32 v1, 0xffff, v1
+; GFX7-SDAG-NEXT: v_mad_u64_u32 v[0:1], s[4:5], v0, v1, v[2:3]
+; GFX7-SDAG-NEXT: s_setpc_b64 s[30:31]
;
-; GFX8-LABEL: mul_u24_known_24bit_add64:
-; GFX8: ; %bb.0:
-; GFX8-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GFX8-NEXT: v_mul_hi_u32_u24_sdwa v4, v0, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_0 src1_sel:WORD_0
-; GFX8-NEXT: v_mul_u32_u24_sdwa v0, v0, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_0 src1_sel:WORD_0
-; GFX8-NEXT: v_add_u32_e32 v0, vcc, v0, v2
-; GFX8-NEXT: v_addc_u32_e32 v1, vcc, v4, v3, vcc
-; GFX8-NEXT: s_setpc_b64 s[30:31]
+; GFX7-GISEL-LABEL: mul_u24_known_24bit_add64:
+; GFX7-GISEL: ; %bb.0:
+; GFX7-GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX7-GISEL-NEXT: v_and_b32_e32 v0, 0xffff, v0
+; GFX7-GISEL-NEXT: v_and_b32_e32 v1, 0xffff, v1
+; GFX7-GISEL-NEXT: v_mul_hi_u32_u24_e32 v4, v0, v1
+; GFX7-GISEL-NEXT: v_mul_u32_u24_e32 v0, v0, v1
+; GFX7-GISEL-NEXT: v_add_i32_e32 v0, vcc, v0, v2
+; GFX7-GISEL-NEXT: v_addc_u32_e32 v1, vcc, v4, v3, vcc
+; GFX7-GISEL-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX8-SDAG-LABEL: mul_u24_known_24bit_add64:
+; GFX8-SDAG: ; %bb.0:
+; GFX8-SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX8-SDAG-NEXT: v_and_b32_e32 v0, 0xffff, v0
+; GFX8-SDAG-NEXT: v_and_b32_e32 v1, 0xffff, v1
+; GFX8-SDAG-NEXT: v_mad_u64_u32 v[0:1], s[4:5], v0, v1, v[2:3]
+; GFX8-SDAG-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX8-GISEL-LABEL: mul_u24_known_24bit_add64:
+; GFX8-GISEL: ; %bb.0:
+; GFX8-GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX8-GISEL-NEXT: v_mul_hi_u32_u24_sdwa v4, v0, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_0 src1_sel:WORD_0
+; GFX8-GISEL-NEXT: v_mul_u32_u24_sdwa v0, v0, v1 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_0 src1_sel:WORD_0
+; GFX8-GISEL-NEXT: v_add_u32_e32 v0, vcc, v0, v2
+; GFX8-GISEL-NEXT: v_addc_u32_e32 v1, vcc, v4, v3, vcc
+; GFX8-GISEL-NEXT: s_setpc_b64 s[30:31]
;
; GFX9-SDAG-LABEL: mul_u24_known_24bit_add64:
; GFX9-SDAG: ; %bb.0:
@@ -12749,7 +12859,7 @@ define i64 @mul_u24_known_24bit_add64(i16 %x.s, i16 %y.s, i64 %z) {
; GFX1250-SDAG-FAKE16-NEXT: v_and_b32_e32 v0, 0xffff, v0
; GFX1250-SDAG-FAKE16-NEXT: v_and_b32_e32 v1, 0xffff, v1
; GFX1250-SDAG-FAKE16-NEXT: s_delay_alu instid0(VALU_DEP_1)
-; GFX1250-SDAG-FAKE16-NEXT: v_mad_co_u64_u32 v[0:1], null, v0, v1, v[2:3]
+; GFX1250-SDAG-FAKE16-NEXT: v_mad_nc_u64_u32 v[0:1], v0, v1, v[2:3]
; GFX1250-SDAG-FAKE16-NEXT: s_set_pc_i64 s[30:31]
;
; GFX1250-GISEL-FAKE16-LABEL: mul_u24_known_24bit_add64:
@@ -12771,7 +12881,7 @@ define i64 @mul_u24_known_24bit_add64(i16 %x.s, i16 %y.s, i64 %z) {
; GFX1250-SDAG-REAL16-NEXT: v_cvt_u32_u16_e32 v0, v0.l
; GFX1250-SDAG-REAL16-NEXT: v_cvt_u32_u16_e32 v1, v1.l
; GFX1250-SDAG-REAL16-NEXT: s_delay_alu instid0(VALU_DEP_1)
-; GFX1250-SDAG-REAL16-NEXT: v_mad_co_u64_u32 v[0:1], null, v0, v1, v[2:3]
+; GFX1250-SDAG-REAL16-NEXT: v_mad_nc_u64_u32 v[0:1], v0, v1, v[2:3]
; GFX1250-SDAG-REAL16-NEXT: s_set_pc_i64 s[30:31]
;
; GFX1250-GISEL-REAL16-LABEL: mul_u24_known_24bit_add64:
@@ -12796,7 +12906,6 @@ declare i64 @llvm.amdgcn.mul.u24.i64(i32, i32)
declare i32 @llvm.amdgcn.mul.u24(i32, i32)
;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:
-; GFX6: {{.*}}
; GFX7: {{.*}}
; GFX900: {{.*}}
; GFX90A: {{.*}}
diff --git a/llvm/test/CodeGen/AMDGPU/mad_u64_u32_mul24.ll b/llvm/test/CodeGen/AMDGPU/mad_u64_u32_mul24.ll
new file mode 100644
index 00000000000000..27d0f7b9f6a8d4
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/mad_u64_u32_mul24.ll
@@ -0,0 +1,126 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc -mtriple=amdgpu9.08 < %s | FileCheck -check-prefix=GCN %s
+
+; performMulCombine rewrites an i64 mul with 24-bit operands into
+; build_pair(mul24, mulhi24). That product feeding an add or ptradd must still
+; fuse into v_mad_u64_u32 / v_mad_i64_i32, not a mul24 + carry chain.
+
+declare i32 @llvm.amdgcn.workitem.id.x()
+
+; (add (build_pair mul_u24, mulhi_u24), z) -> v_mad_u64_u32
+define amdgpu_kernel void @mul24_u_add_to_mad(ptr addrspace(1) %p, ptr addrspace(1) %q) {
+; GCN-LABEL: mul24_u_add_to_mad:
+; GCN: ; %bb.0:
+; GCN-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24
+; GCN-NEXT: v_lshlrev_b32_e32 v2, 3, v0
+; GCN-NEXT: s_waitcnt lgkmcnt(0)
+; GCN-NEXT: global_load_dword v3, v2, s[0:1] offset:4
+; GCN-NEXT: global_load_dwordx2 v[0:1], v2, s[2:3]
+; GCN-NEXT: s_movk_i32 s2, 0xd1
+; GCN-NEXT: s_waitcnt vmcnt(1)
+; GCN-NEXT: v_lshrrev_b32_e32 v3, 8, v3
+; GCN-NEXT: s_waitcnt vmcnt(0)
+; GCN-NEXT: v_mad_u64_u32 v[0:1], s[2:3], v3, s2, v[0:1]
+; GCN-NEXT: global_store_dwordx2 v2, v[0:1], s[0:1]
+; GCN-NEXT: s_endpgm
+ %tid = call i32 @llvm.amdgcn.workitem.id.x()
+ %pg = getelementptr i64, ptr addrspace(1) %p, i32 %tid
+ %qg = getelementptr i64, ptr addrspace(1) %q, i32 %tid
+ %x = load i64, ptr addrspace(1) %pg
+ %acc = load i64, ptr addrspace(1) %qg
+ %hi = lshr i64 %x, 40 ; <= 24 significant bits (u24)
+ %mul = mul i64 %hi, 209 ; 209 fits in 24 bits
+ %add = add i64 %mul, %acc
+ store i64 %add, ptr addrspace(1) %pg
+ ret void
+}
+
+; Negative signed case. Operand is built with a logical shift so bits [31:24]
+; are zero, not sign. MUL_I24 signs from bit 23 but mad from bit 31, so no fuse.
+define amdgpu_kernel void @mul24_i_add_no_mad(ptr addrspace(1) %p, ptr addrspace(1) %q) {
+; GCN-LABEL: mul24_i_add_no_mad:
+; GCN: ; %bb.0:
+; GCN-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24
+; GCN-NEXT: v_lshlrev_b32_e32 v2, 3, v0
+; GCN-NEXT: s_waitcnt lgkmcnt(0)
+; GCN-NEXT: global_load_dword v3, v2, s[0:1] offset:4
+; GCN-NEXT: global_load_dwordx2 v[0:1], v2, s[2:3]
+; GCN-NEXT: s_waitcnt vmcnt(1)
+; GCN-NEXT: v_lshrrev_b32_e32 v3, 8, v3
+; GCN-NEXT: v_mul_hi_i32_i24_e32 v4, 0x64, v3
+; GCN-NEXT: v_mul_i32_i24_e32 v3, 0x64, v3
+; GCN-NEXT: s_waitcnt vmcnt(0)
+; GCN-NEXT: v_add_co_u32_e32 v0, vcc, v3, v0
+; GCN-NEXT: v_addc_co_u32_e32 v1, vcc, v4, v1, vcc
+; GCN-NEXT: global_store_dwordx2 v2, v[0:1], s[0:1]
+; GCN-NEXT: s_endpgm
+ %tid = call i32 @llvm.amdgcn.workitem.id.x()
+ %pg = getelementptr i64, ptr addrspace(1) %p, i32 %tid
+ %qg = getelementptr i64, ptr addrspace(1) %q, i32 %tid
+ %x = load i64, ptr addrspace(1) %pg
+ %acc = load i64, ptr addrspace(1) %qg
+ %sh = ashr i64 %x, 40 ; <= 24 significant signed bits (i24)
+ %mul = mul i64 %sh, 100 ; 100 fits in signed 24 bits
+ %add = add i64 %mul, %acc
+ store i64 %add, ptr addrspace(1) %pg
+ ret void
+}
+
+; ptradd form: (ptradd base, (build_pair mul_u24, mulhi_u24)) -> v_mad_u64_u32
+define amdgpu_kernel void @mul24_u_ptradd_to_mad(ptr addrspace(1) %p) {
+; GCN-LABEL: mul24_u_ptradd_to_mad:
+; GCN: ; %bb.0:
+; GCN-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24
+; GCN-NEXT: v_lshlrev_b32_e32 v0, 3, v0
+; GCN-NEXT: s_movk_i32 s2, 0xd1
+; GCN-NEXT: s_waitcnt lgkmcnt(0)
+; GCN-NEXT: global_load_dwordx2 v[0:1], v0, s[0:1]
+; GCN-NEXT: v_mov_b32_e32 v3, s1
+; GCN-NEXT: v_mov_b32_e32 v2, s0
+; GCN-NEXT: s_waitcnt vmcnt(0)
+; GCN-NEXT: v_lshrrev_b32_e32 v4, 8, v1
+; GCN-NEXT: v_mad_u64_u32 v[2:3], s[0:1], v4, s2, v[2:3]
+; GCN-NEXT: global_store_dwordx2 v[2:3], v[0:1], off
+; GCN-NEXT: s_endpgm
+ %tid = call i32 @llvm.amdgcn.workitem.id.x()
+ %pg = getelementptr i64, ptr addrspace(1) %p, i32 %tid
+ %x = load i64, ptr addrspace(1) %pg
+ %hi = lshr i64 %x, 40
+ %off = mul i64 %hi, 209
+ %gep = getelementptr i8, ptr addrspace(1) %p, i64 %off
+ store i64 %x, ptr addrspace(1) %gep
+ ret void
+}
+
+; Product also used by a store. The add still fuses and the product is
+; recomputed for the store.
+define amdgpu_kernel void @mul24_u_add_extra_use(ptr addrspace(1) %p, ptr addrspace(1) %q) {
+; GCN-LABEL: mul24_u_add_extra_use:
+; GCN: ; %bb.0:
+; GCN-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24
+; GCN-NEXT: v_lshlrev_b32_e32 v4, 3, v0
+; GCN-NEXT: s_movk_i32 s4, 0xd1
+; GCN-NEXT: s_waitcnt lgkmcnt(0)
+; GCN-NEXT: global_load_dword v2, v4, s[0:1] offset:4
+; GCN-NEXT: global_load_dwordx2 v[0:1], v4, s[2:3]
+; GCN-NEXT: s_waitcnt vmcnt(1)
+; GCN-NEXT: v_lshrrev_b32_e32 v2, 8, v2
+; GCN-NEXT: s_waitcnt vmcnt(0)
+; GCN-NEXT: v_mad_u64_u32 v[0:1], s[4:5], v2, s4, v[0:1]
+; GCN-NEXT: v_mul_hi_u32_u24_e32 v3, 0xd1, v2
+; GCN-NEXT: v_mul_u32_u24_e32 v2, 0xd1, v2
+; GCN-NEXT: global_store_dwordx2 v4, v[0:1], s[0:1]
+; GCN-NEXT: global_store_dwordx2 v4, v[2:3], s[2:3]
+; GCN-NEXT: s_endpgm
+ %tid = call i32 @llvm.amdgcn.workitem.id.x()
+ %pg = getelementptr i64, ptr addrspace(1) %p, i32 %tid
+ %qg = getelementptr i64, ptr addrspace(1) %q, i32 %tid
+ %x = load i64, ptr addrspace(1) %pg
+ %acc = load i64, ptr addrspace(1) %qg
+ %hi = lshr i64 %x, 40
+ %mul = mul i64 %hi, 209
+ %add = add i64 %mul, %acc
+ store i64 %add, ptr addrspace(1) %pg
+ store i64 %mul, ptr addrspace(1) %qg ; second use of %mul
+ ret void
+}
More information about the llvm-commits
mailing list