[llvm] [NVPTX] Enhance `mul.wide` and `mad.wide` peepholes (PR #150477)
Justin Fargnoli via llvm-commits
llvm-commits at lists.llvm.org
Mon Jul 28 16:08:26 PDT 2025
https://github.com/justinfargnoli updated https://github.com/llvm/llvm-project/pull/150477
>From 1c1cd4a744cec9b9cb04a26162ddff1330507454 Mon Sep 17 00:00:00 2001
From: Justin Fargnoli <jfargnoli at nvidia.com>
Date: Thu, 24 Jul 2025 17:19:53 +0000
Subject: [PATCH 1/6] Initial commit
---
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 37 ++++++-
llvm/lib/Target/NVPTX/NVPTXISelLowering.h | 2 +
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td | 20 ++++
llvm/test/CodeGen/NVPTX/combine-ext-mad.ll | 117 ++++++++++++++++++++
4 files changed, 171 insertions(+), 5 deletions(-)
create mode 100644 llvm/test/CodeGen/NVPTX/combine-ext-mad.ll
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 7aa06f9079b09..fb9f4c844a1a4 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -1101,6 +1101,8 @@ const char *NVPTXTargetLowering::getTargetNodeName(unsigned Opcode) const {
MAKE_CASE(NVPTXISD::SETP_BF16X2)
MAKE_CASE(NVPTXISD::MUL_WIDE_SIGNED)
MAKE_CASE(NVPTXISD::MUL_WIDE_UNSIGNED)
+ MAKE_CASE(NVPTXISD::MAD_WIDE_UNSIGNED)
+ MAKE_CASE(NVPTXISD::MAD_WIDE_SIGNED)
MAKE_CASE(NVPTXISD::BrxEnd)
MAKE_CASE(NVPTXISD::BrxItem)
MAKE_CASE(NVPTXISD::BrxStart)
@@ -4885,6 +4887,30 @@ static bool isConstZero(const SDValue &Operand) {
return Const && Const->getZExtValue() == 0;
}
+static SDValue
+PerformMADCombineWithOperands(SDNode *N, SDValue N0, SDValue N1,
+ TargetLowering::DAGCombinerInfo &DCI) {
+ assert(N->getOpcode() == ISD::ADD);
+ if (!(N0.getOpcode() == ISD::ZERO_EXTEND ||
+ N0.getOpcode() == ISD::ANY_EXTEND ||
+ N0.getOpcode() == ISD::SIGN_EXTEND))
+ return SDValue();
+ if (N->getValueType(0) != MVT::i64)
+ return SDValue();
+ SDValue M = N0.getOperand(0);
+ if (M.getOpcode() != ISD::MUL)
+ return SDValue();
+ if (M.getValueType() != MVT::i32)
+ return SDValue();
+
+ unsigned Opcode = NVPTXISD::MAD_WIDE_UNSIGNED;
+ if (N0.getOpcode() == ISD::SIGN_EXTEND)
+ Opcode = NVPTXISD::MAD_WIDE_SIGNED;
+ SDValue Mul = N0.getOperand(0);
+ return DCI.DAG.getNode(Opcode, SDLoc(N), N->getValueType(0),
+ Mul.getOperand(0), Mul.getOperand(1), N1);
+}
+
/// PerformADDCombineWithOperands - Try DAG combinations for an ADD with
/// operands N0 and N1. This is a helper for PerformADDCombine that is
/// called with the default operands, and if that fails, with commuted
@@ -4905,6 +4931,9 @@ PerformADDCombineWithOperands(SDNode *N, SDValue N0, SDValue N1,
// -> (select cond, c, (add (mul a, b), c))
//
if (N0.getOpcode() == ISD::SELECT) {
+ // Skip non-integer, non-scalar case
+ if (VT.isVector() || VT != MVT::i32)
+ return SDValue();
unsigned ZeroOpNum;
if (isConstZero(N0->getOperand(1)))
ZeroOpNum = 1;
@@ -4926,6 +4955,9 @@ PerformADDCombineWithOperands(SDNode *N, SDValue N0, SDValue N1,
((ZeroOpNum == 1) ? MAD : N1));
}
+ if (SDValue V = PerformMADCombineWithOperands(N, N0, N1, DCI))
+ return V;
+
return SDValue();
}
@@ -5274,11 +5306,6 @@ static SDValue PerformADDCombine(SDNode *N,
SDValue N0 = N->getOperand(0);
SDValue N1 = N->getOperand(1);
- // Skip non-integer, non-scalar case
- EVT VT = N0.getValueType();
- if (VT.isVector() || VT != MVT::i32)
- return SDValue();
-
// First try with the default operand order.
if (SDValue Result = PerformADDCombineWithOperands(N, N0, N1, DCI))
return Result;
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
index bc3548c0272bb..39c3787641ad0 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
@@ -48,6 +48,8 @@ enum NodeType : unsigned {
FSHR_CLAMP,
MUL_WIDE_SIGNED,
MUL_WIDE_UNSIGNED,
+ MAD_WIDE_UNSIGNED,
+ MAD_WIDE_SIGNED,
SETP_F16X2,
SETP_BF16X2,
BFI,
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index a5bb83dfadb84..4ef650f6d7397 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -990,6 +990,26 @@ defm MAD32 : MAD<"mad.lo.s32", i32, B32, i32imm>;
defm MAD64 : MAD<"mad.lo.s64", i64, B64, i64imm>;
}
+def SDTMadWide : SDTypeProfile<1, 3, [SDTCisInt<0>, SDTCisSameAs<0, 3>, SDTCisInt<1>, SDTCisSameAs<1, 2>]>;
+def mad_wide_signed : SDNode<"NVPTXISD::MAD_WIDE_SIGNED", SDTMadWide>;
+def mad_wide_unsigned : SDNode<"NVPTXISD::MAD_WIDE_UNSIGNED", SDTMadWide>;
+
+multiclass MAD_WIDE<string PtxSuffix, SDNode Op, ValueType BigVT, NVPTXRegClass BigReg, ValueType SmallVT, NVPTXRegClass SmallReg, Operand SmallImm> {
+ def rrr:
+ BasicNVPTXInst<(outs BigReg:$dst),
+ (ins SmallReg:$a, SmallReg:$b, BigReg:$c),
+ "mad.wide." # PtxSuffix,
+ [(set BigVT:$dst, (Op SmallVT:$a, SmallVT:$b, BigVT:$c))]>;
+ def rir:
+ BasicNVPTXInst<(outs BigReg:$dst),
+ (ins SmallReg:$a, SmallImm:$b, BigReg:$c),
+ "mad.wide." # PtxSuffix,
+ [(set BigVT:$dst, (Op SmallVT:$a, imm:$b, BigVT:$c))]>;
+}
+
+defm MAD_WIDE_UNSIGNED_32 : MAD_WIDE<"u32", mad_wide_unsigned, i64, Int64Regs, i32, Int32Regs, i32imm>;
+defm MAD_WIDE_SIGNED_32 : MAD_WIDE<"s32", mad_wide_signed, i64, Int64Regs, i32, Int32Regs, i32imm>;
+
foreach t = [I16RT, I32RT, I64RT] in {
def NEG_S # t.Size :
BasicNVPTXInst<(outs t.RC:$dst), (ins t.RC:$src),
diff --git a/llvm/test/CodeGen/NVPTX/combine-ext-mad.ll b/llvm/test/CodeGen/NVPTX/combine-ext-mad.ll
new file mode 100644
index 0000000000000..c6d656ef65725
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/combine-ext-mad.ll
@@ -0,0 +1,117 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+; RUN: llc < %s | FileCheck %s
+
+target triple = "nvptx64-nvidia-cuda"
+
+define i64 @t1(i32 %a, i32 %b, i64 %c) {
+; CHECK-LABEL: t1(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [t1_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [t1_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [t1_param_2];
+; CHECK-NEXT: mad.wide.s32 %rd2, %r1, %r2, %rd1;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
+; CHECK-NEXT: ret;
+ %mul = mul i32 %a, %b
+ %sext = sext i32 %mul to i64
+ %add = add i64 %c, %sext
+ ret i64 %add
+}
+
+define i64 @t2(i32 %a, i32 %b, i64 %c) {
+; CHECK-LABEL: t2(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [t2_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [t2_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [t2_param_2];
+; CHECK-NEXT: mad.wide.s32 %rd2, %r1, %r2, %rd1;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
+; CHECK-NEXT: ret;
+ %mul = mul i32 %a, %b
+ %sext = sext i32 %mul to i64
+ %add = add i64 %sext, %c
+ ret i64 %add
+}
+
+define i64 @t3(i32 %a, i32 %b) {
+; CHECK-LABEL: t3(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [t3_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [t3_param_1];
+; CHECK-NEXT: mov.b64 %rd1, 1;
+; CHECK-NEXT: mad.wide.s32 %rd2, %r1, %r2, %rd1;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
+; CHECK-NEXT: ret;
+ %mul = mul i32 %a, %b
+ %sext = sext i32 %mul to i64
+ %add = add i64 1, %sext
+ ret i64 %add
+}
+
+define i64 @t4(i32 %a, i64 %c) {
+; CHECK-LABEL: t4(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [t4_param_0];
+; CHECK-NEXT: ld.param.b64 %rd1, [t4_param_1];
+; CHECK-NEXT: mad.wide.s32 %rd2, %r1, 3, %rd1;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
+; CHECK-NEXT: ret;
+ %mul = mul i32 %a, 3
+ %sext = sext i32 %mul to i64
+ %add = add i64 %c, %sext
+ ret i64 %add
+}
+
+define i64 @t5(i32 %a, i32 %b, i64 %c) {
+; CHECK-LABEL: t5(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [t5_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [t5_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [t5_param_2];
+; CHECK-NEXT: mad.wide.u32 %rd2, %r1, %r2, %rd1;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
+; CHECK-NEXT: ret;
+ %mul = mul i32 %a, %b
+ %zext = zext i32 %mul to i64
+ %add = add i64 %c, %zext
+ ret i64 %add
+}
+
+define i64 @t6(i32 %a, i32 %b, i64 %c) {
+; CHECK-LABEL: t6(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [t6_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [t6_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [t6_param_2];
+; CHECK-NEXT: mad.wide.u32 %rd2, %r1, %r2, %rd1;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
+; CHECK-NEXT: ret;
+ %mul = mul i32 %a, %b
+ %zext = zext i32 %mul to i64
+ %add = add i64 %zext, %c
+ ret i64 %add
+}
>From 9612afe0f10aa312e8f35d66a9dd900b1e0384dd Mon Sep 17 00:00:00 2001
From: Justin Fargnoli <jfargnoli at nvidia.com>
Date: Fri, 25 Jul 2025 20:06:07 +0000
Subject: [PATCH 2/6] Switch to (add (mul.wide a, b), c) -> (mad.wide a, b, c)
approach
---
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 90 ++-
llvm/lib/Target/NVPTX/NVPTXISelLowering.h | 2 -
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td | 28 +-
llvm/test/CodeGen/NVPTX/bug26185-2.ll | 9 +-
llvm/test/CodeGen/NVPTX/combine-ext-mad.ll | 117 ----
llvm/test/CodeGen/NVPTX/combine-wide.ll | 693 +++++++++++++++++++
llvm/test/CodeGen/NVPTX/local-stack-frame.ll | 7 +-
llvm/test/CodeGen/NVPTX/vector-loads.ll | 11 +-
8 files changed, 780 insertions(+), 177 deletions(-)
delete mode 100644 llvm/test/CodeGen/NVPTX/combine-ext-mad.ll
create mode 100644 llvm/test/CodeGen/NVPTX/combine-wide.ll
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index fb9f4c844a1a4..e9659e7ffc840 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -841,7 +841,8 @@ NVPTXTargetLowering::NVPTXTargetLowering(const NVPTXTargetMachine &TM,
setTargetDAGCombine({ISD::ADD, ISD::AND, ISD::EXTRACT_VECTOR_ELT, ISD::FADD,
ISD::MUL, ISD::SHL, ISD::SREM, ISD::UREM, ISD::VSELECT,
ISD::BUILD_VECTOR, ISD::ADDRSPACECAST, ISD::LOAD,
- ISD::STORE});
+ ISD::STORE, ISD::ZERO_EXTEND, ISD::SIGN_EXTEND,
+ ISD::ANY_EXTEND});
// setcc for f16x2 and bf16x2 needs special handling to prevent
// legalizer's attempt to scalarize it due to v2i1 not being legal.
@@ -1101,8 +1102,6 @@ const char *NVPTXTargetLowering::getTargetNodeName(unsigned Opcode) const {
MAKE_CASE(NVPTXISD::SETP_BF16X2)
MAKE_CASE(NVPTXISD::MUL_WIDE_SIGNED)
MAKE_CASE(NVPTXISD::MUL_WIDE_UNSIGNED)
- MAKE_CASE(NVPTXISD::MAD_WIDE_UNSIGNED)
- MAKE_CASE(NVPTXISD::MAD_WIDE_SIGNED)
MAKE_CASE(NVPTXISD::BrxEnd)
MAKE_CASE(NVPTXISD::BrxItem)
MAKE_CASE(NVPTXISD::BrxStart)
@@ -4887,30 +4886,6 @@ static bool isConstZero(const SDValue &Operand) {
return Const && Const->getZExtValue() == 0;
}
-static SDValue
-PerformMADCombineWithOperands(SDNode *N, SDValue N0, SDValue N1,
- TargetLowering::DAGCombinerInfo &DCI) {
- assert(N->getOpcode() == ISD::ADD);
- if (!(N0.getOpcode() == ISD::ZERO_EXTEND ||
- N0.getOpcode() == ISD::ANY_EXTEND ||
- N0.getOpcode() == ISD::SIGN_EXTEND))
- return SDValue();
- if (N->getValueType(0) != MVT::i64)
- return SDValue();
- SDValue M = N0.getOperand(0);
- if (M.getOpcode() != ISD::MUL)
- return SDValue();
- if (M.getValueType() != MVT::i32)
- return SDValue();
-
- unsigned Opcode = NVPTXISD::MAD_WIDE_UNSIGNED;
- if (N0.getOpcode() == ISD::SIGN_EXTEND)
- Opcode = NVPTXISD::MAD_WIDE_SIGNED;
- SDValue Mul = N0.getOperand(0);
- return DCI.DAG.getNode(Opcode, SDLoc(N), N->getValueType(0),
- Mul.getOperand(0), Mul.getOperand(1), N1);
-}
-
/// PerformADDCombineWithOperands - Try DAG combinations for an ADD with
/// operands N0 and N1. This is a helper for PerformADDCombine that is
/// called with the default operands, and if that fails, with commuted
@@ -4931,9 +4906,6 @@ PerformADDCombineWithOperands(SDNode *N, SDValue N0, SDValue N1,
// -> (select cond, c, (add (mul a, b), c))
//
if (N0.getOpcode() == ISD::SELECT) {
- // Skip non-integer, non-scalar case
- if (VT.isVector() || VT != MVT::i32)
- return SDValue();
unsigned ZeroOpNum;
if (isConstZero(N0->getOperand(1)))
ZeroOpNum = 1;
@@ -4955,9 +4927,6 @@ PerformADDCombineWithOperands(SDNode *N, SDValue N0, SDValue N1,
((ZeroOpNum == 1) ? MAD : N1));
}
- if (SDValue V = PerformMADCombineWithOperands(N, N0, N1, DCI))
- return V;
-
return SDValue();
}
@@ -5306,6 +5275,10 @@ static SDValue PerformADDCombine(SDNode *N,
SDValue N0 = N->getOperand(0);
SDValue N1 = N->getOperand(1);
+ // Skip non-integer, non-scalar case
+ if (N->getValueType(0).isVector() || N->getValueType(0) != MVT::i32)
+ return SDValue();
+
// First try with the default operand order.
if (SDValue Result = PerformADDCombineWithOperands(N, N0, N1, DCI))
return Result;
@@ -5435,6 +5408,53 @@ static SDValue PerformREMCombine(SDNode *N,
return SDValue();
}
+// (any_extend|sign_extend|zero_extend (mul|shl) x, y) -> (mul.wide x, y)
+static SDValue
+PerformExtendMULWIDECombine(SDNode *N, TargetLowering::DAGCombinerInfo &DCI) {
+ unsigned ExtOpcode = N->getOpcode();
+ assert(ExtOpcode == ISD::ANY_EXTEND || ExtOpcode == ISD::SIGN_EXTEND ||
+ ExtOpcode == ISD::ZERO_EXTEND);
+ EVT ToVT = N->getValueType(0);
+ if (!(ToVT == MVT::i32 || ToVT == MVT::i64))
+ return SDValue();
+ SDValue Op = N->getOperand(0);
+ if (!(Op.getOpcode() == ISD::MUL || Op.getOpcode() == ISD::SHL))
+ return SDValue();
+ if (Op.getOpcode() == ISD::SHL && !isa<ConstantSDNode>(Op.getOperand(1)))
+ return SDValue();
+ EVT FromVT = Op.getValueType();
+ if (!(FromVT == MVT::i16 || FromVT == MVT::i32))
+ return SDValue();
+ if (ExtOpcode == ISD::SIGN_EXTEND && !Op->getFlags().hasNoSignedWrap())
+ return SDValue();
+ if (ExtOpcode == ISD::ZERO_EXTEND && !Op->getFlags().hasNoUnsignedWrap())
+ return SDValue();
+ if (ExtOpcode == ISD::ANY_EXTEND && !Op->getFlags().hasNoSignedWrap() &&
+ !Op->getFlags().hasNoUnsignedWrap())
+ return SDValue();
+
+ SDLoc DL(N);
+ unsigned Opcode = 0;
+ if (ExtOpcode == ISD::SIGN_EXTEND)
+ Opcode = NVPTXISD::MUL_WIDE_SIGNED;
+ else if (ExtOpcode == ISD::ZERO_EXTEND)
+ Opcode = NVPTXISD::MUL_WIDE_UNSIGNED;
+ else if (ExtOpcode == ISD::ANY_EXTEND && Op->getFlags().hasNoUnsignedWrap())
+ Opcode = NVPTXISD::MUL_WIDE_UNSIGNED;
+ else if (ExtOpcode == ISD::ANY_EXTEND && Op->getFlags().hasNoSignedWrap())
+ Opcode = NVPTXISD::MUL_WIDE_SIGNED;
+ else
+ assert(false);
+ SDValue RHS = Op.getOperand(1);
+ if (Op.getOpcode() == ISD::SHL) {
+ const auto ShiftAmt =
+ cast<ConstantSDNode>(Op.getOperand(1))->getZExtValue();
+ const auto MulVal = APInt(ToVT.getSizeInBits(), 1) << ShiftAmt;
+ RHS = DCI.DAG.getConstant(MulVal, DL, ToVT);
+ }
+ return DCI.DAG.getNode(Opcode, DL, ToVT, Op.getOperand(0), RHS);
+}
+
enum OperandSignedness {
Signed = 0,
Unsigned,
@@ -5957,6 +5977,10 @@ SDValue NVPTXTargetLowering::PerformDAGCombine(SDNode *N,
return combineADDRSPACECAST(N, DCI);
case ISD::AND:
return PerformANDCombine(N, DCI);
+ case ISD::ANY_EXTEND:
+ case ISD::SIGN_EXTEND:
+ case ISD::ZERO_EXTEND:
+ return PerformExtendMULWIDECombine(N, DCI);
case ISD::BUILD_VECTOR:
return PerformBUILD_VECTORCombine(N, DCI);
case ISD::EXTRACT_VECTOR_ELT:
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
index 39c3787641ad0..bc3548c0272bb 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
@@ -48,8 +48,6 @@ enum NodeType : unsigned {
FSHR_CLAMP,
MUL_WIDE_SIGNED,
MUL_WIDE_UNSIGNED,
- MAD_WIDE_UNSIGNED,
- MAD_WIDE_SIGNED,
SETP_F16X2,
SETP_BF16X2,
BFI,
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index 4ef650f6d7397..b353ab26afc88 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -859,7 +859,7 @@ def MULWIDEU32Imm :
def MULWIDEU32Imm32 :
BasicNVPTXInst<(outs B32:$dst), (ins B16:$a, i32imm:$b), "mul.wide.u16">;
-def SDTMulWide : SDTypeProfile<1, 2, [SDTCisSameAs<1, 2>]>;
+def SDTMulWide : SDTypeProfile<1, 2, [SDTCisInt<0>, SDTCisInt<1>, SDTCisSameAs<1, 2>]>;
def mul_wide_signed : SDNode<"NVPTXISD::MUL_WIDE_SIGNED", SDTMulWide>;
def mul_wide_unsigned : SDNode<"NVPTXISD::MUL_WIDE_UNSIGNED", SDTMulWide>;
@@ -990,25 +990,33 @@ defm MAD32 : MAD<"mad.lo.s32", i32, B32, i32imm>;
defm MAD64 : MAD<"mad.lo.s64", i64, B64, i64imm>;
}
-def SDTMadWide : SDTypeProfile<1, 3, [SDTCisInt<0>, SDTCisSameAs<0, 3>, SDTCisInt<1>, SDTCisSameAs<1, 2>]>;
-def mad_wide_signed : SDNode<"NVPTXISD::MAD_WIDE_SIGNED", SDTMadWide>;
-def mad_wide_unsigned : SDNode<"NVPTXISD::MAD_WIDE_UNSIGNED", SDTMadWide>;
-
-multiclass MAD_WIDE<string PtxSuffix, SDNode Op, ValueType BigVT, NVPTXRegClass BigReg, ValueType SmallVT, NVPTXRegClass SmallReg, Operand SmallImm> {
+multiclass MAD_WIDE<string PtxSuffix, SDNode Op, ValueType BigVT, NVPTXRegClass BigReg, Operand BigImm, ValueType SmallVT, NVPTXRegClass SmallReg, Operand SmallImm> {
def rrr:
BasicNVPTXInst<(outs BigReg:$dst),
(ins SmallReg:$a, SmallReg:$b, BigReg:$c),
"mad.wide." # PtxSuffix,
- [(set BigVT:$dst, (Op SmallVT:$a, SmallVT:$b, BigVT:$c))]>;
+ [(set BigVT:$dst, (add (Op SmallVT:$a, SmallVT:$b), BigVT:$c))]>;
+ def rri:
+ BasicNVPTXInst<(outs BigReg:$dst),
+ (ins SmallReg:$a, SmallReg:$b, BigImm:$c),
+ "mad.wide." # PtxSuffix,
+ [(set BigVT:$dst, (add (Op SmallVT:$a, SmallVT:$b), imm:$c))]>;
def rir:
BasicNVPTXInst<(outs BigReg:$dst),
(ins SmallReg:$a, SmallImm:$b, BigReg:$c),
"mad.wide." # PtxSuffix,
- [(set BigVT:$dst, (Op SmallVT:$a, imm:$b, BigVT:$c))]>;
+ [(set BigVT:$dst, (add (Op SmallVT:$a, imm:$b), BigVT:$c))]>;
+ def rii:
+ BasicNVPTXInst<(outs BigReg:$dst),
+ (ins SmallReg:$a, SmallImm:$b, BigImm:$c),
+ "mad.wide." # PtxSuffix,
+ [(set BigVT:$dst, (add (Op SmallVT:$a, imm:$b), imm:$c))]>;
}
-defm MAD_WIDE_UNSIGNED_32 : MAD_WIDE<"u32", mad_wide_unsigned, i64, Int64Regs, i32, Int32Regs, i32imm>;
-defm MAD_WIDE_SIGNED_32 : MAD_WIDE<"s32", mad_wide_signed, i64, Int64Regs, i32, Int32Regs, i32imm>;
+defm MAD_WIDE_U16 : MAD_WIDE<"u16", mul_wide_unsigned, i32, B32, i32imm, i16, B16, i16imm>;
+defm MAD_WIDE_S16 : MAD_WIDE<"s16", mul_wide_signed, i32, B32, i32imm, i16, B16, i16imm>;
+defm MAD_WIDE_U32 : MAD_WIDE<"u32", mul_wide_unsigned, i64, B64, i64imm, i32, B32, i32imm>;
+defm MAD_WIDE_S32 : MAD_WIDE<"s32", mul_wide_signed, i64, B64, i64imm, i32, B32, i32imm>;
foreach t = [I16RT, I32RT, I64RT] in {
def NEG_S # t.Size :
diff --git a/llvm/test/CodeGen/NVPTX/bug26185-2.ll b/llvm/test/CodeGen/NVPTX/bug26185-2.ll
index 4e11f58f85ee0..46172b1af1236 100644
--- a/llvm/test/CodeGen/NVPTX/bug26185-2.ll
+++ b/llvm/test/CodeGen/NVPTX/bug26185-2.ll
@@ -16,7 +16,7 @@ define ptx_kernel void @spam(ptr addrspace(1) noalias nocapture readonly %arg, p
; CHECK: .maxntid 1, 1, 1
; CHECK-NEXT: {
; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .b64 %rd<9>;
+; CHECK-NEXT: .reg .b64 %rd<8>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0: // %bb
; CHECK-NEXT: ld.param.b64 %rd1, [spam_param_0];
@@ -25,10 +25,9 @@ define ptx_kernel void @spam(ptr addrspace(1) noalias nocapture readonly %arg, p
; CHECK-NEXT: add.s64 %rd4, %rd1, %rd3;
; CHECK-NEXT: ld.param.b64 %rd5, [spam_param_1];
; CHECK-NEXT: ld.global.nc.s16 %r1, [%rd4+16];
-; CHECK-NEXT: mul.wide.s32 %rd6, %r1, %r1;
-; CHECK-NEXT: ld.global.b64 %rd7, [%rd5];
-; CHECK-NEXT: add.s64 %rd8, %rd6, %rd7;
-; CHECK-NEXT: st.global.b64 [%rd5], %rd8;
+; CHECK-NEXT: ld.global.b64 %rd6, [%rd5];
+; CHECK-NEXT: mad.wide.s32 %rd7, %r1, %r1, %rd6;
+; CHECK-NEXT: st.global.b64 [%rd5], %rd7;
; CHECK-NEXT: ret;
bb:
%tmp5 = add nsw i64 %arg3, 8
diff --git a/llvm/test/CodeGen/NVPTX/combine-ext-mad.ll b/llvm/test/CodeGen/NVPTX/combine-ext-mad.ll
deleted file mode 100644
index c6d656ef65725..0000000000000
--- a/llvm/test/CodeGen/NVPTX/combine-ext-mad.ll
+++ /dev/null
@@ -1,117 +0,0 @@
-; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
-; RUN: llc < %s | FileCheck %s
-
-target triple = "nvptx64-nvidia-cuda"
-
-define i64 @t1(i32 %a, i32 %b, i64 %c) {
-; CHECK-LABEL: t1(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .b64 %rd<3>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r1, [t1_param_0];
-; CHECK-NEXT: ld.param.b32 %r2, [t1_param_1];
-; CHECK-NEXT: ld.param.b64 %rd1, [t1_param_2];
-; CHECK-NEXT: mad.wide.s32 %rd2, %r1, %r2, %rd1;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
-; CHECK-NEXT: ret;
- %mul = mul i32 %a, %b
- %sext = sext i32 %mul to i64
- %add = add i64 %c, %sext
- ret i64 %add
-}
-
-define i64 @t2(i32 %a, i32 %b, i64 %c) {
-; CHECK-LABEL: t2(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .b64 %rd<3>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r1, [t2_param_0];
-; CHECK-NEXT: ld.param.b32 %r2, [t2_param_1];
-; CHECK-NEXT: ld.param.b64 %rd1, [t2_param_2];
-; CHECK-NEXT: mad.wide.s32 %rd2, %r1, %r2, %rd1;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
-; CHECK-NEXT: ret;
- %mul = mul i32 %a, %b
- %sext = sext i32 %mul to i64
- %add = add i64 %sext, %c
- ret i64 %add
-}
-
-define i64 @t3(i32 %a, i32 %b) {
-; CHECK-LABEL: t3(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .b64 %rd<3>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r1, [t3_param_0];
-; CHECK-NEXT: ld.param.b32 %r2, [t3_param_1];
-; CHECK-NEXT: mov.b64 %rd1, 1;
-; CHECK-NEXT: mad.wide.s32 %rd2, %r1, %r2, %rd1;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
-; CHECK-NEXT: ret;
- %mul = mul i32 %a, %b
- %sext = sext i32 %mul to i64
- %add = add i64 1, %sext
- ret i64 %add
-}
-
-define i64 @t4(i32 %a, i64 %c) {
-; CHECK-LABEL: t4(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .b64 %rd<3>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r1, [t4_param_0];
-; CHECK-NEXT: ld.param.b64 %rd1, [t4_param_1];
-; CHECK-NEXT: mad.wide.s32 %rd2, %r1, 3, %rd1;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
-; CHECK-NEXT: ret;
- %mul = mul i32 %a, 3
- %sext = sext i32 %mul to i64
- %add = add i64 %c, %sext
- ret i64 %add
-}
-
-define i64 @t5(i32 %a, i32 %b, i64 %c) {
-; CHECK-LABEL: t5(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .b64 %rd<3>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r1, [t5_param_0];
-; CHECK-NEXT: ld.param.b32 %r2, [t5_param_1];
-; CHECK-NEXT: ld.param.b64 %rd1, [t5_param_2];
-; CHECK-NEXT: mad.wide.u32 %rd2, %r1, %r2, %rd1;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
-; CHECK-NEXT: ret;
- %mul = mul i32 %a, %b
- %zext = zext i32 %mul to i64
- %add = add i64 %c, %zext
- ret i64 %add
-}
-
-define i64 @t6(i32 %a, i32 %b, i64 %c) {
-; CHECK-LABEL: t6(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .b64 %rd<3>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r1, [t6_param_0];
-; CHECK-NEXT: ld.param.b32 %r2, [t6_param_1];
-; CHECK-NEXT: ld.param.b64 %rd1, [t6_param_2];
-; CHECK-NEXT: mad.wide.u32 %rd2, %r1, %r2, %rd1;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
-; CHECK-NEXT: ret;
- %mul = mul i32 %a, %b
- %zext = zext i32 %mul to i64
- %add = add i64 %zext, %c
- ret i64 %add
-}
diff --git a/llvm/test/CodeGen/NVPTX/combine-wide.ll b/llvm/test/CodeGen/NVPTX/combine-wide.ll
new file mode 100644
index 0000000000000..218ac4ed567f1
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/combine-wide.ll
@@ -0,0 +1,693 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+; RUN: llc < %s | FileCheck %s
+
+target triple = "nvptx64-nvidia-cuda"
+
+define i64 @t1(i32 %a, i32 %b, i64 %c) {
+; CHECK-LABEL: t1(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [t1_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [t1_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [t1_param_2];
+; CHECK-NEXT: mad.wide.s32 %rd2, %r1, %r2, %rd1;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
+; CHECK-NEXT: ret;
+ %mul = mul nsw i32 %a, %b
+ %sext = sext i32 %mul to i64
+ %add = add i64 %c, %sext
+ ret i64 %add
+}
+
+define i64 @t2(i32 %a, i32 %b, i64 %c) {
+; CHECK-LABEL: t2(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [t2_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [t2_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [t2_param_2];
+; CHECK-NEXT: mad.wide.s32 %rd2, %r1, %r2, %rd1;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
+; CHECK-NEXT: ret;
+ %mul = mul nsw i32 %a, %b
+ %sext = sext i32 %mul to i64
+ %add = add i64 %sext, %c
+ ret i64 %add
+}
+
+define i64 @t3(i32 %a, i32 %b) {
+; CHECK-LABEL: t3(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [t3_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [t3_param_1];
+; CHECK-NEXT: mad.wide.s32 %rd1, %r1, %r2, 1;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
+; CHECK-NEXT: ret;
+ %mul = mul nsw i32 %a, %b
+ %sext = sext i32 %mul to i64
+ %add = add i64 1, %sext
+ ret i64 %add
+}
+
+define i64 @t4(i32 %a, i64 %c) {
+; CHECK-LABEL: t4(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [t4_param_0];
+; CHECK-NEXT: ld.param.b64 %rd1, [t4_param_1];
+; CHECK-NEXT: mad.wide.s32 %rd2, %r1, 3, %rd1;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
+; CHECK-NEXT: ret;
+ %mul = mul nsw i32 %a, 3
+ %sext = sext i32 %mul to i64
+ %add = add i64 %c, %sext
+ ret i64 %add
+}
+
+define i64 @t4_1(i32 %a, i64 %c) {
+; CHECK-LABEL: t4_1(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [t4_1_param_0];
+; CHECK-NEXT: mad.wide.s32 %rd1, %r1, 3, 5;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
+; CHECK-NEXT: ret;
+ %mul = mul nsw i32 %a, 3
+ %sext = sext i32 %mul to i64
+ %add = add i64 5, %sext
+ ret i64 %add
+}
+
+define i64 @t5(i32 %a, i32 %b, i64 %c) {
+; CHECK-LABEL: t5(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [t5_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [t5_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [t5_param_2];
+; CHECK-NEXT: mad.wide.u32 %rd2, %r1, %r2, %rd1;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
+; CHECK-NEXT: ret;
+ %mul = mul nuw i32 %a, %b
+ %zext = zext i32 %mul to i64
+ %add = add i64 %c, %zext
+ ret i64 %add
+}
+
+define i64 @t6(i32 %a, i32 %b, i64 %c) {
+; CHECK-LABEL: t6(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [t6_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [t6_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [t6_param_2];
+; CHECK-NEXT: mad.wide.u32 %rd2, %r1, %r2, %rd1;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
+; CHECK-NEXT: ret;
+ %mul = mul nuw i32 %a, %b
+ %zext = zext i32 %mul to i64
+ %add = add i64 %zext, %c
+ ret i64 %add
+}
+
+define i32 @t7(i16 %a, i16 %b) {
+; CHECK-LABEL: t7(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<4>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [t7_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [t7_param_1];
+; CHECK-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %mul = mul i16 %a, %b
+ %zext = zext i16 %mul to i32
+ ret i32 %zext
+}
+
+define i32 @t8(i16 %a, i16 %b) {
+; CHECK-LABEL: t8(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<4>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [t8_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [t8_param_1];
+; CHECK-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; CHECK-NEXT: cvt.s32.s16 %r1, %rs3;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %mul = mul i16 %a, %b
+ %sext = sext i16 %mul to i32
+ ret i32 %sext
+}
+
+define i64 @t9(i32 %a, i32 %b) {
+; CHECK-LABEL: t9(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [t9_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [t9_param_1];
+; CHECK-NEXT: mul.lo.s32 %r3, %r1, %r2;
+; CHECK-NEXT: cvt.u64.u32 %rd1, %r3;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
+; CHECK-NEXT: ret;
+ %mul = mul i32 %a, %b
+ %zext = zext i32 %mul to i64
+ ret i64 %zext
+}
+
+define i64 @t10(i32 %a, i32 %b) {
+; CHECK-LABEL: t10(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [t10_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [t10_param_1];
+; CHECK-NEXT: mul.lo.s32 %r3, %r1, %r2;
+; CHECK-NEXT: cvt.s64.s32 %rd1, %r3;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
+; CHECK-NEXT: ret;
+ %mul = mul i32 %a, %b
+ %sext = sext i32 %mul to i64
+ ret i64 %sext
+}
+
+define i32 @t11(i16 %a, i16 %b) {
+; CHECK-LABEL: t11(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<4>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [t11_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [t11_param_1];
+; CHECK-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %mul = mul nsw i16 %a, %b
+ %zext = zext i16 %mul to i32
+ ret i32 %zext
+}
+
+define i32 @t12(i16 %a, i16 %b) {
+; CHECK-LABEL: t12(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [t12_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [t12_param_1];
+; CHECK-NEXT: mul.wide.s16 %r1, %rs1, %rs2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %mul = mul nsw i16 %a, %b
+ %sext = sext i16 %mul to i32
+ ret i32 %sext
+}
+
+define i64 @t13(i32 %a, i32 %b) {
+; CHECK-LABEL: t13(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [t13_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [t13_param_1];
+; CHECK-NEXT: mul.lo.s32 %r3, %r1, %r2;
+; CHECK-NEXT: cvt.u64.u32 %rd1, %r3;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
+; CHECK-NEXT: ret;
+ %mul = mul nsw i32 %a, %b
+ %zext = zext i32 %mul to i64
+ ret i64 %zext
+}
+
+define i64 @t14(i32 %a, i32 %b) {
+; CHECK-LABEL: t14(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [t14_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [t14_param_1];
+; CHECK-NEXT: mul.wide.s32 %rd1, %r1, %r2;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
+; CHECK-NEXT: ret;
+ %mul = mul nsw i32 %a, %b
+ %sext = sext i32 %mul to i64
+ ret i64 %sext
+}
+
+define i32 @t15(i16 %a, i16 %b) {
+; CHECK-LABEL: t15(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [t15_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [t15_param_1];
+; CHECK-NEXT: mul.wide.u16 %r1, %rs1, %rs2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %mul = mul nuw i16 %a, %b
+ %zext = zext i16 %mul to i32
+ ret i32 %zext
+}
+
+define i32 @t16(i16 %a, i16 %b) {
+; CHECK-LABEL: t16(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<4>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [t16_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [t16_param_1];
+; CHECK-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; CHECK-NEXT: cvt.s32.s16 %r1, %rs3;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %mul = mul nuw i16 %a, %b
+ %sext = sext i16 %mul to i32
+ ret i32 %sext
+}
+
+define i64 @t17(i32 %a, i32 %b) {
+; CHECK-LABEL: t17(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [t17_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [t17_param_1];
+; CHECK-NEXT: mul.wide.u32 %rd1, %r1, %r2;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
+; CHECK-NEXT: ret;
+ %mul = mul nuw i32 %a, %b
+ %zext = zext i32 %mul to i64
+ ret i64 %zext
+}
+
+define i64 @t18(i32 %a, i32 %b) {
+; CHECK-LABEL: t18(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [t18_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [t18_param_1];
+; CHECK-NEXT: mul.lo.s32 %r3, %r1, %r2;
+; CHECK-NEXT: cvt.s64.s32 %rd1, %r3;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
+; CHECK-NEXT: ret;
+ %mul = mul nuw i32 %a, %b
+ %sext = sext i32 %mul to i64
+ ret i64 %sext
+}
+
+define i32 @t19(i16 %a, i16 %b) {
+; CHECK-LABEL: t19(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<4>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [t19_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [t19_param_1];
+; CHECK-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %mul = mul i16 %a, %b
+ %zext = zext i16 %mul to i32
+ ret i32 %zext
+}
+
+define i32 @t20(i16 %a) {
+; CHECK-LABEL: t20(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [t20_param_0];
+; CHECK-NEXT: shl.b16 %rs2, %rs1, 4;
+; CHECK-NEXT: cvt.s32.s16 %r1, %rs2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %mul = shl i16 %a, 4
+ %sext = sext i16 %mul to i32
+ ret i32 %sext
+}
+
+define i64 @t21(i32 %a) {
+; CHECK-LABEL: t21(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [t21_param_0];
+; CHECK-NEXT: shl.b32 %r2, %r1, 4;
+; CHECK-NEXT: cvt.u64.u32 %rd1, %r2;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
+; CHECK-NEXT: ret;
+ %mul = shl i32 %a, 4
+ %zext = zext i32 %mul to i64
+ ret i64 %zext
+}
+
+define i64 @t22(i32 %a) {
+; CHECK-LABEL: t22(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [t22_param_0];
+; CHECK-NEXT: shl.b32 %r2, %r1, 4;
+; CHECK-NEXT: cvt.s64.s32 %rd1, %r2;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
+; CHECK-NEXT: ret;
+ %mul = shl i32 %a, 4
+ %sext = sext i32 %mul to i64
+ ret i64 %sext
+}
+
+define i32 @t23(i16 %a, i16 %b) {
+; CHECK-LABEL: t23(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [t23_param_0];
+; CHECK-NEXT: shl.b16 %rs2, %rs1, 4;
+; CHECK-NEXT: cvt.u32.u16 %r1, %rs2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %mul = shl nsw i16 %a, 4
+ %zext = zext i16 %mul to i32
+ ret i32 %zext
+}
+
+define i32 @t24(i16 %a, i16 %b) {
+; CHECK-LABEL: t24(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [t24_param_0];
+; CHECK-NEXT: mul.wide.s16 %r1, %rs1, 16;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %mul = shl nsw i16 %a, 4
+ %sext = sext i16 %mul to i32
+ ret i32 %sext
+}
+
+define i64 @t25(i32 %a) {
+; CHECK-LABEL: t25(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [t25_param_0];
+; CHECK-NEXT: shl.b32 %r2, %r1, 4;
+; CHECK-NEXT: cvt.u64.u32 %rd1, %r2;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
+; CHECK-NEXT: ret;
+ %mul = shl nsw i32 %a, 4
+ %zext = zext i32 %mul to i64
+ ret i64 %zext
+}
+
+define i64 @t26(i32 %a) {
+; CHECK-LABEL: t26(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [t26_param_0];
+; CHECK-NEXT: mul.wide.s32 %rd1, %r1, 16;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
+; CHECK-NEXT: ret;
+ %mul = shl nsw i32 %a, 4
+ %sext = sext i32 %mul to i64
+ ret i64 %sext
+}
+
+define i32 @t27(i16 %a, i16 %b) {
+; CHECK-LABEL: t27(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [t27_param_0];
+; CHECK-NEXT: mul.wide.u16 %r1, %rs1, 16;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %mul = shl nuw i16 %a, 4
+ %zext = zext i16 %mul to i32
+ ret i32 %zext
+}
+
+define i32 @t28(i16 %a, i16 %b) {
+; CHECK-LABEL: t28(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [t28_param_0];
+; CHECK-NEXT: shl.b16 %rs2, %rs1, 4;
+; CHECK-NEXT: cvt.s32.s16 %r1, %rs2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %mul = shl nuw i16 %a, 4
+ %sext = sext i16 %mul to i32
+ ret i32 %sext
+}
+
+define i64 @t29(i32 %a) {
+; CHECK-LABEL: t29(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [t29_param_0];
+; CHECK-NEXT: mul.wide.u32 %rd1, %r1, 16;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
+; CHECK-NEXT: ret;
+ %mul = shl nuw i32 %a, 4
+ %zext = zext i32 %mul to i64
+ ret i64 %zext
+}
+
+define i64 @t30(i32 %a) {
+; CHECK-LABEL: t30(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [t30_param_0];
+; CHECK-NEXT: shl.b32 %r2, %r1, 4;
+; CHECK-NEXT: cvt.s64.s32 %rd1, %r2;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
+; CHECK-NEXT: ret;
+ %mul = shl nuw i32 %a, 4
+ %sext = sext i32 %mul to i64
+ ret i64 %sext
+}
+
+define i64 @t31(i32 %a, i32 %b) {
+; CHECK-LABEL: t31(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [t31_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [t31_param_1];
+; CHECK-NEXT: shl.b32 %r3, %r1, %r2;
+; CHECK-NEXT: cvt.s64.s32 %rd1, %r3;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
+; CHECK-NEXT: ret;
+ %mul = shl nuw i32 %a, %b
+ %sext = sext i32 %mul to i64
+ ret i64 %sext
+}
+
+define i32 @t32(i16 %a, i16 %b, i32 %c) {
+; CHECK-LABEL: t32(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [t32_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [t32_param_1];
+; CHECK-NEXT: ld.param.b32 %r1, [t32_param_2];
+; CHECK-NEXT: mad.wide.s16 %r2, %rs1, %rs2, %r1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %mul = mul nsw i16 %a, %b
+ %sext = sext i16 %mul to i32
+ %add = add i32 %c, %sext
+ ret i32 %add
+}
+
+define i32 @t33(i16 %a, i16 %b, i32 %c) {
+; CHECK-LABEL: t33(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [t33_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [t33_param_1];
+; CHECK-NEXT: ld.param.b32 %r1, [t33_param_2];
+; CHECK-NEXT: mad.wide.s16 %r2, %rs1, %rs2, %r1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %mul = mul nsw i16 %a, %b
+ %sext = sext i16 %mul to i32
+ %add = add i32 %c, %sext
+ ret i32 %add
+}
+
+define i32 @t34(i16 %a, i16 %b) {
+; CHECK-LABEL: t34(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [t34_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [t34_param_1];
+; CHECK-NEXT: mad.wide.s16 %r1, %rs1, %rs2, 1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %mul = mul nsw i16 %a, %b
+ %sext = sext i16 %mul to i32
+ %add = add i32 1, %sext
+ ret i32 %add
+}
+
+define i32 @t35(i16 %a, i32 %c) {
+; CHECK-LABEL: t35(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [t35_param_0];
+; CHECK-NEXT: ld.param.b32 %r1, [t35_param_1];
+; CHECK-NEXT: mad.wide.s16 %r2, %rs1, 3, %r1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %mul = mul nsw i16 %a, 3
+ %sext = sext i16 %mul to i32
+ %add = add i32 %c, %sext
+ ret i32 %add
+}
+
+define i32 @t36(i16 %a, i32 %c) {
+; CHECK-LABEL: t36(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [t36_param_0];
+; CHECK-NEXT: mad.wide.s16 %r1, %rs1, 3, 5;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %mul = mul nsw i16 %a, 3
+ %sext = sext i16 %mul to i32
+ %add = add i32 5, %sext
+ ret i32 %add
+}
+
+define i32 @t37(i16 %a, i16 %b, i32 %c) {
+; CHECK-LABEL: t37(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [t37_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [t37_param_1];
+; CHECK-NEXT: ld.param.b32 %r1, [t37_param_2];
+; CHECK-NEXT: mad.wide.u16 %r2, %rs1, %rs2, %r1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %mul = mul nuw i16 %a, %b
+ %zext = zext i16 %mul to i32
+ %add = add i32 %c, %zext
+ ret i32 %add
+}
+
+define i32 @t38(i16 %a, i16 %b, i32 %c) {
+; CHECK-LABEL: t38(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b16 %rs1, [t38_param_0];
+; CHECK-NEXT: ld.param.b16 %rs2, [t38_param_1];
+; CHECK-NEXT: ld.param.b32 %r1, [t38_param_2];
+; CHECK-NEXT: mad.wide.u16 %r2, %rs1, %rs2, %r1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %mul = mul nuw i16 %a, %b
+ %zext = zext i16 %mul to i32
+ %add = add i32 %zext, %c
+ ret i32 %add
+}
diff --git a/llvm/test/CodeGen/NVPTX/local-stack-frame.ll b/llvm/test/CodeGen/NVPTX/local-stack-frame.ll
index 5c3017310d0a3..ae069cf956c36 100644
--- a/llvm/test/CodeGen/NVPTX/local-stack-frame.ll
+++ b/llvm/test/CodeGen/NVPTX/local-stack-frame.ll
@@ -114,15 +114,14 @@ define void @foo3(i32 %a) {
; PTX64-NEXT: .reg .b64 %SP;
; PTX64-NEXT: .reg .b64 %SPL;
; PTX64-NEXT: .reg .b32 %r<2>;
-; PTX64-NEXT: .reg .b64 %rd<5>;
+; PTX64-NEXT: .reg .b64 %rd<4>;
; PTX64-EMPTY:
; PTX64-NEXT: // %bb.0:
; PTX64-NEXT: mov.b64 %SPL, __local_depot2;
; PTX64-NEXT: ld.param.b32 %r1, [foo3_param_0];
; PTX64-NEXT: add.u64 %rd2, %SPL, 0;
-; PTX64-NEXT: mul.wide.s32 %rd3, %r1, 4;
-; PTX64-NEXT: add.s64 %rd4, %rd2, %rd3;
-; PTX64-NEXT: st.local.b32 [%rd4], %r1;
+; PTX64-NEXT: mad.wide.s32 %rd3, %r1, 4, %rd2;
+; PTX64-NEXT: st.local.b32 [%rd3], %r1;
; PTX64-NEXT: ret;
%local = alloca [3 x i32], align 4
%1 = getelementptr inbounds i32, ptr %local, i32 %a
diff --git a/llvm/test/CodeGen/NVPTX/vector-loads.ll b/llvm/test/CodeGen/NVPTX/vector-loads.ll
index e16fc74325416..6f0dff78d5569 100644
--- a/llvm/test/CodeGen/NVPTX/vector-loads.ll
+++ b/llvm/test/CodeGen/NVPTX/vector-loads.ll
@@ -154,7 +154,7 @@ define void @foo_complex(ptr nocapture readonly align 16 dereferenceable(1342177
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<4>;
; CHECK-NEXT: .reg .b32 %r<8>;
-; CHECK-NEXT: .reg .b64 %rd<6>;
+; CHECK-NEXT: .reg .b64 %rd<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b64 %rd1, [foo_complex_param_0];
@@ -166,12 +166,11 @@ define void @foo_complex(ptr nocapture readonly align 16 dereferenceable(1342177
; CHECK-NEXT: shl.b32 %r6, %r1, 1;
; CHECK-NEXT: or.b32 %r7, %r5, %r6;
; CHECK-NEXT: cvt.u64.u32 %rd2, %r7;
-; CHECK-NEXT: mul.wide.u32 %rd3, %r3, 131072;
-; CHECK-NEXT: add.s64 %rd4, %rd1, %rd3;
-; CHECK-NEXT: add.s64 %rd5, %rd4, %rd2;
-; CHECK-NEXT: ld.v2.b8 {%rs1, %rs2}, [%rd5+128];
+; CHECK-NEXT: mad.wide.u32 %rd3, %r3, 131072, %rd1;
+; CHECK-NEXT: add.s64 %rd4, %rd3, %rd2;
+; CHECK-NEXT: ld.v2.b8 {%rs1, %rs2}, [%rd4+128];
; CHECK-NEXT: max.u16 %rs3, %rs1, %rs2;
-; CHECK-NEXT: st.b8 [%rd5+129], %rs3;
+; CHECK-NEXT: st.b8 [%rd4+129], %rs3;
; CHECK-NEXT: ret;
%t0 = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x(), !range !1
%t1 = tail call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
>From 13d6a1172b7e2421289d506c6a7f31856e7d614d Mon Sep 17 00:00:00 2001
From: Justin Fargnoli <jfargnoli at nvidia.com>
Date: Fri, 25 Jul 2025 20:07:16 +0000
Subject: [PATCH 3/6] Original output of combine-wide.ll
---
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 1 +
llvm/test/CodeGen/NVPTX/combine-wide.ll | 196 ++++++++++++--------
2 files changed, 117 insertions(+), 80 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index e9659e7ffc840..16dfa160c5368 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -5411,6 +5411,7 @@ static SDValue PerformREMCombine(SDNode *N,
// (any_extend|sign_extend|zero_extend (mul|shl) x, y) -> (mul.wide x, y)
static SDValue
PerformExtendMULWIDECombine(SDNode *N, TargetLowering::DAGCombinerInfo &DCI) {
+ return SDValue();
unsigned ExtOpcode = N->getOpcode();
assert(ExtOpcode == ISD::ANY_EXTEND || ExtOpcode == ISD::SIGN_EXTEND ||
ExtOpcode == ISD::ZERO_EXTEND);
diff --git a/llvm/test/CodeGen/NVPTX/combine-wide.ll b/llvm/test/CodeGen/NVPTX/combine-wide.ll
index 218ac4ed567f1..f6da9bede66e8 100644
--- a/llvm/test/CodeGen/NVPTX/combine-wide.ll
+++ b/llvm/test/CodeGen/NVPTX/combine-wide.ll
@@ -6,15 +6,17 @@ target triple = "nvptx64-nvidia-cuda"
define i64 @t1(i32 %a, i32 %b, i64 %c) {
; CHECK-LABEL: t1(
; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-NEXT: .reg .b64 %rd<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [t1_param_0];
; CHECK-NEXT: ld.param.b32 %r2, [t1_param_1];
-; CHECK-NEXT: ld.param.b64 %rd1, [t1_param_2];
-; CHECK-NEXT: mad.wide.s32 %rd2, %r1, %r2, %rd1;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
+; CHECK-NEXT: mul.lo.s32 %r3, %r1, %r2;
+; CHECK-NEXT: cvt.s64.s32 %rd1, %r3;
+; CHECK-NEXT: ld.param.b64 %rd2, [t1_param_2];
+; CHECK-NEXT: add.s64 %rd3, %rd2, %rd1;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
; CHECK-NEXT: ret;
%mul = mul nsw i32 %a, %b
%sext = sext i32 %mul to i64
@@ -25,15 +27,17 @@ define i64 @t1(i32 %a, i32 %b, i64 %c) {
define i64 @t2(i32 %a, i32 %b, i64 %c) {
; CHECK-LABEL: t2(
; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-NEXT: .reg .b64 %rd<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [t2_param_0];
; CHECK-NEXT: ld.param.b32 %r2, [t2_param_1];
-; CHECK-NEXT: ld.param.b64 %rd1, [t2_param_2];
-; CHECK-NEXT: mad.wide.s32 %rd2, %r1, %r2, %rd1;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
+; CHECK-NEXT: mul.lo.s32 %r3, %r1, %r2;
+; CHECK-NEXT: cvt.s64.s32 %rd1, %r3;
+; CHECK-NEXT: ld.param.b64 %rd2, [t2_param_2];
+; CHECK-NEXT: add.s64 %rd3, %rd1, %rd2;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
; CHECK-NEXT: ret;
%mul = mul nsw i32 %a, %b
%sext = sext i32 %mul to i64
@@ -44,14 +48,16 @@ define i64 @t2(i32 %a, i32 %b, i64 %c) {
define i64 @t3(i32 %a, i32 %b) {
; CHECK-LABEL: t3(
; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [t3_param_0];
; CHECK-NEXT: ld.param.b32 %r2, [t3_param_1];
-; CHECK-NEXT: mad.wide.s32 %rd1, %r1, %r2, 1;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
+; CHECK-NEXT: mul.lo.s32 %r3, %r1, %r2;
+; CHECK-NEXT: cvt.s64.s32 %rd1, %r3;
+; CHECK-NEXT: add.s64 %rd2, %rd1, 1;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
; CHECK-NEXT: ret;
%mul = mul nsw i32 %a, %b
%sext = sext i32 %mul to i64
@@ -62,14 +68,16 @@ define i64 @t3(i32 %a, i32 %b) {
define i64 @t4(i32 %a, i64 %c) {
; CHECK-LABEL: t4(
; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [t4_param_0];
+; CHECK-NEXT: mul.lo.s32 %r2, %r1, 3;
; CHECK-NEXT: ld.param.b64 %rd1, [t4_param_1];
-; CHECK-NEXT: mad.wide.s32 %rd2, %r1, 3, %rd1;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
+; CHECK-NEXT: cvt.s64.s32 %rd2, %r2;
+; CHECK-NEXT: add.s64 %rd3, %rd1, %rd2;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
; CHECK-NEXT: ret;
%mul = mul nsw i32 %a, 3
%sext = sext i32 %mul to i64
@@ -80,13 +88,15 @@ define i64 @t4(i32 %a, i64 %c) {
define i64 @t4_1(i32 %a, i64 %c) {
; CHECK-LABEL: t4_1(
; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [t4_1_param_0];
-; CHECK-NEXT: mad.wide.s32 %rd1, %r1, 3, 5;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
+; CHECK-NEXT: mul.lo.s32 %r2, %r1, 3;
+; CHECK-NEXT: cvt.s64.s32 %rd1, %r2;
+; CHECK-NEXT: add.s64 %rd2, %rd1, 5;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
; CHECK-NEXT: ret;
%mul = mul nsw i32 %a, 3
%sext = sext i32 %mul to i64
@@ -97,15 +107,17 @@ define i64 @t4_1(i32 %a, i64 %c) {
define i64 @t5(i32 %a, i32 %b, i64 %c) {
; CHECK-LABEL: t5(
; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-NEXT: .reg .b64 %rd<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [t5_param_0];
; CHECK-NEXT: ld.param.b32 %r2, [t5_param_1];
-; CHECK-NEXT: ld.param.b64 %rd1, [t5_param_2];
-; CHECK-NEXT: mad.wide.u32 %rd2, %r1, %r2, %rd1;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
+; CHECK-NEXT: mul.lo.s32 %r3, %r1, %r2;
+; CHECK-NEXT: cvt.u64.u32 %rd1, %r3;
+; CHECK-NEXT: ld.param.b64 %rd2, [t5_param_2];
+; CHECK-NEXT: add.s64 %rd3, %rd2, %rd1;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
; CHECK-NEXT: ret;
%mul = mul nuw i32 %a, %b
%zext = zext i32 %mul to i64
@@ -116,15 +128,17 @@ define i64 @t5(i32 %a, i32 %b, i64 %c) {
define i64 @t6(i32 %a, i32 %b, i64 %c) {
; CHECK-LABEL: t6(
; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-NEXT: .reg .b64 %rd<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [t6_param_0];
; CHECK-NEXT: ld.param.b32 %r2, [t6_param_1];
-; CHECK-NEXT: ld.param.b64 %rd1, [t6_param_2];
-; CHECK-NEXT: mad.wide.u32 %rd2, %r1, %r2, %rd1;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
+; CHECK-NEXT: mul.lo.s32 %r3, %r1, %r2;
+; CHECK-NEXT: cvt.u64.u32 %rd1, %r3;
+; CHECK-NEXT: ld.param.b64 %rd2, [t6_param_2];
+; CHECK-NEXT: add.s64 %rd3, %rd1, %rd2;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
; CHECK-NEXT: ret;
%mul = mul nuw i32 %a, %b
%zext = zext i32 %mul to i64
@@ -225,13 +239,14 @@ define i32 @t11(i16 %a, i16 %b) {
define i32 @t12(i16 %a, i16 %b) {
; CHECK-LABEL: t12(
; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b16 %rs<4>;
; CHECK-NEXT: .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [t12_param_0];
; CHECK-NEXT: ld.param.b16 %rs2, [t12_param_1];
-; CHECK-NEXT: mul.wide.s16 %r1, %rs1, %rs2;
+; CHECK-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; CHECK-NEXT: cvt.s32.s16 %r1, %rs3;
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
; CHECK-NEXT: ret;
%mul = mul nsw i16 %a, %b
@@ -260,13 +275,14 @@ define i64 @t13(i32 %a, i32 %b) {
define i64 @t14(i32 %a, i32 %b) {
; CHECK-LABEL: t14(
; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b32 %r<4>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [t14_param_0];
; CHECK-NEXT: ld.param.b32 %r2, [t14_param_1];
-; CHECK-NEXT: mul.wide.s32 %rd1, %r1, %r2;
+; CHECK-NEXT: mul.lo.s32 %r3, %r1, %r2;
+; CHECK-NEXT: cvt.s64.s32 %rd1, %r3;
; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
; CHECK-NEXT: ret;
%mul = mul nsw i32 %a, %b
@@ -277,13 +293,14 @@ define i64 @t14(i32 %a, i32 %b) {
define i32 @t15(i16 %a, i16 %b) {
; CHECK-LABEL: t15(
; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b16 %rs<4>;
; CHECK-NEXT: .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [t15_param_0];
; CHECK-NEXT: ld.param.b16 %rs2, [t15_param_1];
-; CHECK-NEXT: mul.wide.u16 %r1, %rs1, %rs2;
+; CHECK-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
; CHECK-NEXT: ret;
%mul = mul nuw i16 %a, %b
@@ -312,13 +329,14 @@ define i32 @t16(i16 %a, i16 %b) {
define i64 @t17(i32 %a, i32 %b) {
; CHECK-LABEL: t17(
; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b32 %r<4>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [t17_param_0];
; CHECK-NEXT: ld.param.b32 %r2, [t17_param_1];
-; CHECK-NEXT: mul.wide.u32 %rd1, %r1, %r2;
+; CHECK-NEXT: mul.lo.s32 %r3, %r1, %r2;
+; CHECK-NEXT: cvt.u64.u32 %rd1, %r3;
; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
; CHECK-NEXT: ret;
%mul = mul nuw i32 %a, %b
@@ -433,12 +451,13 @@ define i32 @t23(i16 %a, i16 %b) {
define i32 @t24(i16 %a, i16 %b) {
; CHECK-LABEL: t24(
; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b16 %rs<3>;
; CHECK-NEXT: .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [t24_param_0];
-; CHECK-NEXT: mul.wide.s16 %r1, %rs1, 16;
+; CHECK-NEXT: shl.b16 %rs2, %rs1, 4;
+; CHECK-NEXT: cvt.s32.s16 %r1, %rs2;
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
; CHECK-NEXT: ret;
%mul = shl nsw i16 %a, 4
@@ -466,12 +485,13 @@ define i64 @t25(i32 %a) {
define i64 @t26(i32 %a) {
; CHECK-LABEL: t26(
; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [t26_param_0];
-; CHECK-NEXT: mul.wide.s32 %rd1, %r1, 16;
+; CHECK-NEXT: shl.b32 %r2, %r1, 4;
+; CHECK-NEXT: cvt.s64.s32 %rd1, %r2;
; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
; CHECK-NEXT: ret;
%mul = shl nsw i32 %a, 4
@@ -482,12 +502,13 @@ define i64 @t26(i32 %a) {
define i32 @t27(i16 %a, i16 %b) {
; CHECK-LABEL: t27(
; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b16 %rs<3>;
; CHECK-NEXT: .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [t27_param_0];
-; CHECK-NEXT: mul.wide.u16 %r1, %rs1, 16;
+; CHECK-NEXT: shl.b16 %rs2, %rs1, 4;
+; CHECK-NEXT: cvt.u32.u16 %r1, %rs2;
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
; CHECK-NEXT: ret;
%mul = shl nuw i16 %a, 4
@@ -515,12 +536,13 @@ define i32 @t28(i16 %a, i16 %b) {
define i64 @t29(i32 %a) {
; CHECK-LABEL: t29(
; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [t29_param_0];
-; CHECK-NEXT: mul.wide.u32 %rd1, %r1, 16;
+; CHECK-NEXT: shl.b32 %r2, %r1, 4;
+; CHECK-NEXT: cvt.u64.u32 %rd1, %r2;
; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
; CHECK-NEXT: ret;
%mul = shl nuw i32 %a, 4
@@ -566,15 +588,17 @@ define i64 @t31(i32 %a, i32 %b) {
define i32 @t32(i16 %a, i16 %b, i32 %c) {
; CHECK-LABEL: t32(
; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<3>;
-; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b16 %rs<4>;
+; CHECK-NEXT: .reg .b32 %r<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [t32_param_0];
; CHECK-NEXT: ld.param.b16 %rs2, [t32_param_1];
-; CHECK-NEXT: ld.param.b32 %r1, [t32_param_2];
-; CHECK-NEXT: mad.wide.s16 %r2, %rs1, %rs2, %r1;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; CHECK-NEXT: cvt.s32.s16 %r1, %rs3;
+; CHECK-NEXT: ld.param.b32 %r2, [t32_param_2];
+; CHECK-NEXT: add.s32 %r3, %r2, %r1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
; CHECK-NEXT: ret;
%mul = mul nsw i16 %a, %b
%sext = sext i16 %mul to i32
@@ -585,15 +609,17 @@ define i32 @t32(i16 %a, i16 %b, i32 %c) {
define i32 @t33(i16 %a, i16 %b, i32 %c) {
; CHECK-LABEL: t33(
; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<3>;
-; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b16 %rs<4>;
+; CHECK-NEXT: .reg .b32 %r<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [t33_param_0];
; CHECK-NEXT: ld.param.b16 %rs2, [t33_param_1];
-; CHECK-NEXT: ld.param.b32 %r1, [t33_param_2];
-; CHECK-NEXT: mad.wide.s16 %r2, %rs1, %rs2, %r1;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; CHECK-NEXT: cvt.s32.s16 %r1, %rs3;
+; CHECK-NEXT: ld.param.b32 %r2, [t33_param_2];
+; CHECK-NEXT: add.s32 %r3, %r2, %r1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
; CHECK-NEXT: ret;
%mul = mul nsw i16 %a, %b
%sext = sext i16 %mul to i32
@@ -604,14 +630,16 @@ define i32 @t33(i16 %a, i16 %b, i32 %c) {
define i32 @t34(i16 %a, i16 %b) {
; CHECK-LABEL: t34(
; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<3>;
-; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b16 %rs<4>;
+; CHECK-NEXT: .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [t34_param_0];
; CHECK-NEXT: ld.param.b16 %rs2, [t34_param_1];
-; CHECK-NEXT: mad.wide.s16 %r1, %rs1, %rs2, 1;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; CHECK-NEXT: cvt.s32.s16 %r1, %rs3;
+; CHECK-NEXT: add.s32 %r2, %r1, 1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
; CHECK-NEXT: ret;
%mul = mul nsw i16 %a, %b
%sext = sext i16 %mul to i32
@@ -622,14 +650,16 @@ define i32 @t34(i16 %a, i16 %b) {
define i32 @t35(i16 %a, i32 %c) {
; CHECK-LABEL: t35(
; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<2>;
-; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [t35_param_0];
+; CHECK-NEXT: mul.lo.s16 %rs2, %rs1, 3;
; CHECK-NEXT: ld.param.b32 %r1, [t35_param_1];
-; CHECK-NEXT: mad.wide.s16 %r2, %rs1, 3, %r1;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-NEXT: cvt.s32.s16 %r2, %rs2;
+; CHECK-NEXT: add.s32 %r3, %r1, %r2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
; CHECK-NEXT: ret;
%mul = mul nsw i16 %a, 3
%sext = sext i16 %mul to i32
@@ -640,13 +670,15 @@ define i32 @t35(i16 %a, i32 %c) {
define i32 @t36(i16 %a, i32 %c) {
; CHECK-LABEL: t36(
; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<2>;
-; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [t36_param_0];
-; CHECK-NEXT: mad.wide.s16 %r1, %rs1, 3, 5;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: mul.lo.s16 %rs2, %rs1, 3;
+; CHECK-NEXT: cvt.s32.s16 %r1, %rs2;
+; CHECK-NEXT: add.s32 %r2, %r1, 5;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
; CHECK-NEXT: ret;
%mul = mul nsw i16 %a, 3
%sext = sext i16 %mul to i32
@@ -657,15 +689,17 @@ define i32 @t36(i16 %a, i32 %c) {
define i32 @t37(i16 %a, i16 %b, i32 %c) {
; CHECK-LABEL: t37(
; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<3>;
-; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b16 %rs<4>;
+; CHECK-NEXT: .reg .b32 %r<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [t37_param_0];
; CHECK-NEXT: ld.param.b16 %rs2, [t37_param_1];
-; CHECK-NEXT: ld.param.b32 %r1, [t37_param_2];
-; CHECK-NEXT: mad.wide.u16 %r2, %rs1, %rs2, %r1;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;
+; CHECK-NEXT: ld.param.b32 %r2, [t37_param_2];
+; CHECK-NEXT: add.s32 %r3, %r2, %r1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
; CHECK-NEXT: ret;
%mul = mul nuw i16 %a, %b
%zext = zext i16 %mul to i32
@@ -676,15 +710,17 @@ define i32 @t37(i16 %a, i16 %b, i32 %c) {
define i32 @t38(i16 %a, i16 %b, i32 %c) {
; CHECK-LABEL: t38(
; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<3>;
-; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b16 %rs<4>;
+; CHECK-NEXT: .reg .b32 %r<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [t38_param_0];
; CHECK-NEXT: ld.param.b16 %rs2, [t38_param_1];
-; CHECK-NEXT: ld.param.b32 %r1, [t38_param_2];
-; CHECK-NEXT: mad.wide.u16 %r2, %rs1, %rs2, %r1;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;
+; CHECK-NEXT: ld.param.b32 %r2, [t38_param_2];
+; CHECK-NEXT: add.s32 %r3, %r1, %r2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
; CHECK-NEXT: ret;
%mul = mul nuw i16 %a, %b
%zext = zext i16 %mul to i32
>From a69e982ac014a5a957c722f1fea62eb26c7dedeb Mon Sep 17 00:00:00 2001
From: Justin Fargnoli <jfargnoli at nvidia.com>
Date: Fri, 25 Jul 2025 20:07:27 +0000
Subject: [PATCH 4/6] Revert "Original output of combine-wide.ll"
This reverts commit 13d6a1172b7e2421289d506c6a7f31856e7d614d.
---
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 1 -
llvm/test/CodeGen/NVPTX/combine-wide.ll | 196 ++++++++------------
2 files changed, 80 insertions(+), 117 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 16dfa160c5368..e9659e7ffc840 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -5411,7 +5411,6 @@ static SDValue PerformREMCombine(SDNode *N,
// (any_extend|sign_extend|zero_extend (mul|shl) x, y) -> (mul.wide x, y)
static SDValue
PerformExtendMULWIDECombine(SDNode *N, TargetLowering::DAGCombinerInfo &DCI) {
- return SDValue();
unsigned ExtOpcode = N->getOpcode();
assert(ExtOpcode == ISD::ANY_EXTEND || ExtOpcode == ISD::SIGN_EXTEND ||
ExtOpcode == ISD::ZERO_EXTEND);
diff --git a/llvm/test/CodeGen/NVPTX/combine-wide.ll b/llvm/test/CodeGen/NVPTX/combine-wide.ll
index f6da9bede66e8..218ac4ed567f1 100644
--- a/llvm/test/CodeGen/NVPTX/combine-wide.ll
+++ b/llvm/test/CodeGen/NVPTX/combine-wide.ll
@@ -6,17 +6,15 @@ target triple = "nvptx64-nvidia-cuda"
define i64 @t1(i32 %a, i32 %b, i64 %c) {
; CHECK-LABEL: t1(
; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<4>;
-; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [t1_param_0];
; CHECK-NEXT: ld.param.b32 %r2, [t1_param_1];
-; CHECK-NEXT: mul.lo.s32 %r3, %r1, %r2;
-; CHECK-NEXT: cvt.s64.s32 %rd1, %r3;
-; CHECK-NEXT: ld.param.b64 %rd2, [t1_param_2];
-; CHECK-NEXT: add.s64 %rd3, %rd2, %rd1;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
+; CHECK-NEXT: ld.param.b64 %rd1, [t1_param_2];
+; CHECK-NEXT: mad.wide.s32 %rd2, %r1, %r2, %rd1;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
; CHECK-NEXT: ret;
%mul = mul nsw i32 %a, %b
%sext = sext i32 %mul to i64
@@ -27,17 +25,15 @@ define i64 @t1(i32 %a, i32 %b, i64 %c) {
define i64 @t2(i32 %a, i32 %b, i64 %c) {
; CHECK-LABEL: t2(
; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<4>;
-; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [t2_param_0];
; CHECK-NEXT: ld.param.b32 %r2, [t2_param_1];
-; CHECK-NEXT: mul.lo.s32 %r3, %r1, %r2;
-; CHECK-NEXT: cvt.s64.s32 %rd1, %r3;
-; CHECK-NEXT: ld.param.b64 %rd2, [t2_param_2];
-; CHECK-NEXT: add.s64 %rd3, %rd1, %rd2;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
+; CHECK-NEXT: ld.param.b64 %rd1, [t2_param_2];
+; CHECK-NEXT: mad.wide.s32 %rd2, %r1, %r2, %rd1;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
; CHECK-NEXT: ret;
%mul = mul nsw i32 %a, %b
%sext = sext i32 %mul to i64
@@ -48,16 +44,14 @@ define i64 @t2(i32 %a, i32 %b, i64 %c) {
define i64 @t3(i32 %a, i32 %b) {
; CHECK-LABEL: t3(
; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<4>;
-; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [t3_param_0];
; CHECK-NEXT: ld.param.b32 %r2, [t3_param_1];
-; CHECK-NEXT: mul.lo.s32 %r3, %r1, %r2;
-; CHECK-NEXT: cvt.s64.s32 %rd1, %r3;
-; CHECK-NEXT: add.s64 %rd2, %rd1, 1;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
+; CHECK-NEXT: mad.wide.s32 %rd1, %r1, %r2, 1;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
; CHECK-NEXT: ret;
%mul = mul nsw i32 %a, %b
%sext = sext i32 %mul to i64
@@ -68,16 +62,14 @@ define i64 @t3(i32 %a, i32 %b) {
define i64 @t4(i32 %a, i64 %c) {
; CHECK-LABEL: t4(
; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [t4_param_0];
-; CHECK-NEXT: mul.lo.s32 %r2, %r1, 3;
; CHECK-NEXT: ld.param.b64 %rd1, [t4_param_1];
-; CHECK-NEXT: cvt.s64.s32 %rd2, %r2;
-; CHECK-NEXT: add.s64 %rd3, %rd1, %rd2;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
+; CHECK-NEXT: mad.wide.s32 %rd2, %r1, 3, %rd1;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
; CHECK-NEXT: ret;
%mul = mul nsw i32 %a, 3
%sext = sext i32 %mul to i64
@@ -88,15 +80,13 @@ define i64 @t4(i32 %a, i64 %c) {
define i64 @t4_1(i32 %a, i64 %c) {
; CHECK-LABEL: t4_1(
; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [t4_1_param_0];
-; CHECK-NEXT: mul.lo.s32 %r2, %r1, 3;
-; CHECK-NEXT: cvt.s64.s32 %rd1, %r2;
-; CHECK-NEXT: add.s64 %rd2, %rd1, 5;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
+; CHECK-NEXT: mad.wide.s32 %rd1, %r1, 3, 5;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
; CHECK-NEXT: ret;
%mul = mul nsw i32 %a, 3
%sext = sext i32 %mul to i64
@@ -107,17 +97,15 @@ define i64 @t4_1(i32 %a, i64 %c) {
define i64 @t5(i32 %a, i32 %b, i64 %c) {
; CHECK-LABEL: t5(
; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<4>;
-; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [t5_param_0];
; CHECK-NEXT: ld.param.b32 %r2, [t5_param_1];
-; CHECK-NEXT: mul.lo.s32 %r3, %r1, %r2;
-; CHECK-NEXT: cvt.u64.u32 %rd1, %r3;
-; CHECK-NEXT: ld.param.b64 %rd2, [t5_param_2];
-; CHECK-NEXT: add.s64 %rd3, %rd2, %rd1;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
+; CHECK-NEXT: ld.param.b64 %rd1, [t5_param_2];
+; CHECK-NEXT: mad.wide.u32 %rd2, %r1, %r2, %rd1;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
; CHECK-NEXT: ret;
%mul = mul nuw i32 %a, %b
%zext = zext i32 %mul to i64
@@ -128,17 +116,15 @@ define i64 @t5(i32 %a, i32 %b, i64 %c) {
define i64 @t6(i32 %a, i32 %b, i64 %c) {
; CHECK-LABEL: t6(
; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<4>;
-; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [t6_param_0];
; CHECK-NEXT: ld.param.b32 %r2, [t6_param_1];
-; CHECK-NEXT: mul.lo.s32 %r3, %r1, %r2;
-; CHECK-NEXT: cvt.u64.u32 %rd1, %r3;
-; CHECK-NEXT: ld.param.b64 %rd2, [t6_param_2];
-; CHECK-NEXT: add.s64 %rd3, %rd1, %rd2;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
+; CHECK-NEXT: ld.param.b64 %rd1, [t6_param_2];
+; CHECK-NEXT: mad.wide.u32 %rd2, %r1, %r2, %rd1;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
; CHECK-NEXT: ret;
%mul = mul nuw i32 %a, %b
%zext = zext i32 %mul to i64
@@ -239,14 +225,13 @@ define i32 @t11(i16 %a, i16 %b) {
define i32 @t12(i16 %a, i16 %b) {
; CHECK-LABEL: t12(
; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<4>;
+; CHECK-NEXT: .reg .b16 %rs<3>;
; CHECK-NEXT: .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [t12_param_0];
; CHECK-NEXT: ld.param.b16 %rs2, [t12_param_1];
-; CHECK-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
-; CHECK-NEXT: cvt.s32.s16 %r1, %rs3;
+; CHECK-NEXT: mul.wide.s16 %r1, %rs1, %rs2;
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
; CHECK-NEXT: ret;
%mul = mul nsw i16 %a, %b
@@ -275,14 +260,13 @@ define i64 @t13(i32 %a, i32 %b) {
define i64 @t14(i32 %a, i32 %b) {
; CHECK-LABEL: t14(
; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-NEXT: .reg .b32 %r<3>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [t14_param_0];
; CHECK-NEXT: ld.param.b32 %r2, [t14_param_1];
-; CHECK-NEXT: mul.lo.s32 %r3, %r1, %r2;
-; CHECK-NEXT: cvt.s64.s32 %rd1, %r3;
+; CHECK-NEXT: mul.wide.s32 %rd1, %r1, %r2;
; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
; CHECK-NEXT: ret;
%mul = mul nsw i32 %a, %b
@@ -293,14 +277,13 @@ define i64 @t14(i32 %a, i32 %b) {
define i32 @t15(i16 %a, i16 %b) {
; CHECK-LABEL: t15(
; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<4>;
+; CHECK-NEXT: .reg .b16 %rs<3>;
; CHECK-NEXT: .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [t15_param_0];
; CHECK-NEXT: ld.param.b16 %rs2, [t15_param_1];
-; CHECK-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
-; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;
+; CHECK-NEXT: mul.wide.u16 %r1, %rs1, %rs2;
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
; CHECK-NEXT: ret;
%mul = mul nuw i16 %a, %b
@@ -329,14 +312,13 @@ define i32 @t16(i16 %a, i16 %b) {
define i64 @t17(i32 %a, i32 %b) {
; CHECK-LABEL: t17(
; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-NEXT: .reg .b32 %r<3>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [t17_param_0];
; CHECK-NEXT: ld.param.b32 %r2, [t17_param_1];
-; CHECK-NEXT: mul.lo.s32 %r3, %r1, %r2;
-; CHECK-NEXT: cvt.u64.u32 %rd1, %r3;
+; CHECK-NEXT: mul.wide.u32 %rd1, %r1, %r2;
; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
; CHECK-NEXT: ret;
%mul = mul nuw i32 %a, %b
@@ -451,13 +433,12 @@ define i32 @t23(i16 %a, i16 %b) {
define i32 @t24(i16 %a, i16 %b) {
; CHECK-LABEL: t24(
; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b16 %rs<2>;
; CHECK-NEXT: .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [t24_param_0];
-; CHECK-NEXT: shl.b16 %rs2, %rs1, 4;
-; CHECK-NEXT: cvt.s32.s16 %r1, %rs2;
+; CHECK-NEXT: mul.wide.s16 %r1, %rs1, 16;
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
; CHECK-NEXT: ret;
%mul = shl nsw i16 %a, 4
@@ -485,13 +466,12 @@ define i64 @t25(i32 %a) {
define i64 @t26(i32 %a) {
; CHECK-LABEL: t26(
; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b32 %r<2>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [t26_param_0];
-; CHECK-NEXT: shl.b32 %r2, %r1, 4;
-; CHECK-NEXT: cvt.s64.s32 %rd1, %r2;
+; CHECK-NEXT: mul.wide.s32 %rd1, %r1, 16;
; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
; CHECK-NEXT: ret;
%mul = shl nsw i32 %a, 4
@@ -502,13 +482,12 @@ define i64 @t26(i32 %a) {
define i32 @t27(i16 %a, i16 %b) {
; CHECK-LABEL: t27(
; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b16 %rs<2>;
; CHECK-NEXT: .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [t27_param_0];
-; CHECK-NEXT: shl.b16 %rs2, %rs1, 4;
-; CHECK-NEXT: cvt.u32.u16 %r1, %rs2;
+; CHECK-NEXT: mul.wide.u16 %r1, %rs1, 16;
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
; CHECK-NEXT: ret;
%mul = shl nuw i16 %a, 4
@@ -536,13 +515,12 @@ define i32 @t28(i16 %a, i16 %b) {
define i64 @t29(i32 %a) {
; CHECK-LABEL: t29(
; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b32 %r<2>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [t29_param_0];
-; CHECK-NEXT: shl.b32 %r2, %r1, 4;
-; CHECK-NEXT: cvt.u64.u32 %rd1, %r2;
+; CHECK-NEXT: mul.wide.u32 %rd1, %r1, 16;
; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
; CHECK-NEXT: ret;
%mul = shl nuw i32 %a, 4
@@ -588,17 +566,15 @@ define i64 @t31(i32 %a, i32 %b) {
define i32 @t32(i16 %a, i16 %b, i32 %c) {
; CHECK-LABEL: t32(
; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<4>;
-; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [t32_param_0];
; CHECK-NEXT: ld.param.b16 %rs2, [t32_param_1];
-; CHECK-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
-; CHECK-NEXT: cvt.s32.s16 %r1, %rs3;
-; CHECK-NEXT: ld.param.b32 %r2, [t32_param_2];
-; CHECK-NEXT: add.s32 %r3, %r2, %r1;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ld.param.b32 %r1, [t32_param_2];
+; CHECK-NEXT: mad.wide.s16 %r2, %rs1, %rs2, %r1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
; CHECK-NEXT: ret;
%mul = mul nsw i16 %a, %b
%sext = sext i16 %mul to i32
@@ -609,17 +585,15 @@ define i32 @t32(i16 %a, i16 %b, i32 %c) {
define i32 @t33(i16 %a, i16 %b, i32 %c) {
; CHECK-LABEL: t33(
; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<4>;
-; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [t33_param_0];
; CHECK-NEXT: ld.param.b16 %rs2, [t33_param_1];
-; CHECK-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
-; CHECK-NEXT: cvt.s32.s16 %r1, %rs3;
-; CHECK-NEXT: ld.param.b32 %r2, [t33_param_2];
-; CHECK-NEXT: add.s32 %r3, %r2, %r1;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ld.param.b32 %r1, [t33_param_2];
+; CHECK-NEXT: mad.wide.s16 %r2, %rs1, %rs2, %r1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
; CHECK-NEXT: ret;
%mul = mul nsw i16 %a, %b
%sext = sext i16 %mul to i32
@@ -630,16 +604,14 @@ define i32 @t33(i16 %a, i16 %b, i32 %c) {
define i32 @t34(i16 %a, i16 %b) {
; CHECK-LABEL: t34(
; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<4>;
-; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [t34_param_0];
; CHECK-NEXT: ld.param.b16 %rs2, [t34_param_1];
-; CHECK-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
-; CHECK-NEXT: cvt.s32.s16 %r1, %rs3;
-; CHECK-NEXT: add.s32 %r2, %r1, 1;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-NEXT: mad.wide.s16 %r1, %rs1, %rs2, 1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
; CHECK-NEXT: ret;
%mul = mul nsw i16 %a, %b
%sext = sext i16 %mul to i32
@@ -650,16 +622,14 @@ define i32 @t34(i16 %a, i16 %b) {
define i32 @t35(i16 %a, i32 %c) {
; CHECK-LABEL: t35(
; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<3>;
-; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [t35_param_0];
-; CHECK-NEXT: mul.lo.s16 %rs2, %rs1, 3;
; CHECK-NEXT: ld.param.b32 %r1, [t35_param_1];
-; CHECK-NEXT: cvt.s32.s16 %r2, %rs2;
-; CHECK-NEXT: add.s32 %r3, %r1, %r2;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: mad.wide.s16 %r2, %rs1, 3, %r1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
; CHECK-NEXT: ret;
%mul = mul nsw i16 %a, 3
%sext = sext i16 %mul to i32
@@ -670,15 +640,13 @@ define i32 @t35(i16 %a, i32 %c) {
define i32 @t36(i16 %a, i32 %c) {
; CHECK-LABEL: t36(
; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<3>;
-; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [t36_param_0];
-; CHECK-NEXT: mul.lo.s16 %rs2, %rs1, 3;
-; CHECK-NEXT: cvt.s32.s16 %r1, %rs2;
-; CHECK-NEXT: add.s32 %r2, %r1, 5;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-NEXT: mad.wide.s16 %r1, %rs1, 3, 5;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
; CHECK-NEXT: ret;
%mul = mul nsw i16 %a, 3
%sext = sext i16 %mul to i32
@@ -689,17 +657,15 @@ define i32 @t36(i16 %a, i32 %c) {
define i32 @t37(i16 %a, i16 %b, i32 %c) {
; CHECK-LABEL: t37(
; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<4>;
-; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [t37_param_0];
; CHECK-NEXT: ld.param.b16 %rs2, [t37_param_1];
-; CHECK-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
-; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;
-; CHECK-NEXT: ld.param.b32 %r2, [t37_param_2];
-; CHECK-NEXT: add.s32 %r3, %r2, %r1;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ld.param.b32 %r1, [t37_param_2];
+; CHECK-NEXT: mad.wide.u16 %r2, %rs1, %rs2, %r1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
; CHECK-NEXT: ret;
%mul = mul nuw i16 %a, %b
%zext = zext i16 %mul to i32
@@ -710,17 +676,15 @@ define i32 @t37(i16 %a, i16 %b, i32 %c) {
define i32 @t38(i16 %a, i16 %b, i32 %c) {
; CHECK-LABEL: t38(
; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<4>;
-; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [t38_param_0];
; CHECK-NEXT: ld.param.b16 %rs2, [t38_param_1];
-; CHECK-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
-; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;
-; CHECK-NEXT: ld.param.b32 %r2, [t38_param_2];
-; CHECK-NEXT: add.s32 %r3, %r1, %r2;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ld.param.b32 %r1, [t38_param_2];
+; CHECK-NEXT: mad.wide.u16 %r2, %rs1, %rs2, %r1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
; CHECK-NEXT: ret;
%mul = mul nuw i16 %a, %b
%zext = zext i16 %mul to i32
>From 34974eee63e615a8d951064fbd128389e5f97cfb Mon Sep 17 00:00:00 2001
From: Justin Fargnoli <jfargnoli at nvidia.com>
Date: Mon, 28 Jul 2025 23:06:12 +0000
Subject: [PATCH 5/6] Incorporate review feedback
---
llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp | 4 +-
llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.h | 3 -
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 47 +-
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td | 126 +-
llvm/test/CodeGen/NVPTX/combine-wide.ll | 1376 ++++++++++++++-----
5 files changed, 1055 insertions(+), 501 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
index 65e7c56774547..2707cc44f6e11 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
@@ -56,9 +56,7 @@ INITIALIZE_PASS(NVPTXDAGToDAGISelLegacy, DEBUG_TYPE, PASS_NAME, false, false)
NVPTXDAGToDAGISel::NVPTXDAGToDAGISel(NVPTXTargetMachine &tm,
CodeGenOptLevel OptLevel)
- : SelectionDAGISel(tm, OptLevel), TM(tm) {
- doMulWide = (OptLevel > CodeGenOptLevel::None);
-}
+ : SelectionDAGISel(tm, OptLevel), TM(tm) {}
bool NVPTXDAGToDAGISel::runOnMachineFunction(MachineFunction &MF) {
Subtarget = &MF.getSubtarget<NVPTXSubtarget>();
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.h b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.h
index b99b4ef2d3076..ca3a624a1ac4e 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.h
+++ b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.h
@@ -40,9 +40,6 @@ struct NVPTXScopes {
class LLVM_LIBRARY_VISIBILITY NVPTXDAGToDAGISel : public SelectionDAGISel {
const NVPTXTargetMachine &TM;
- // If true, generate mul.wide from sext and mul
- bool doMulWide;
-
NVPTX::DivPrecisionLevel getDivF32Level(const SDNode *N) const;
bool usePrecSqrtF32(const SDNode *N) const;
bool useF32FTZ() const;
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index e9659e7ffc840..b846e7258d10b 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -5276,7 +5276,8 @@ static SDValue PerformADDCombine(SDNode *N,
SDValue N1 = N->getOperand(1);
// Skip non-integer, non-scalar case
- if (N->getValueType(0).isVector() || N->getValueType(0) != MVT::i32)
+ EVT VT = N0.getValueType();
+ if (VT.isVector() || VT != MVT::i32)
return SDValue();
// First try with the default operand order.
@@ -5409,46 +5410,35 @@ static SDValue PerformREMCombine(SDNode *N,
}
// (any_extend|sign_extend|zero_extend (mul|shl) x, y) -> (mul.wide x, y)
-static SDValue
-PerformExtendMULWIDECombine(SDNode *N, TargetLowering::DAGCombinerInfo &DCI) {
- unsigned ExtOpcode = N->getOpcode();
- assert(ExtOpcode == ISD::ANY_EXTEND || ExtOpcode == ISD::SIGN_EXTEND ||
- ExtOpcode == ISD::ZERO_EXTEND);
- EVT ToVT = N->getValueType(0);
- if (!(ToVT == MVT::i32 || ToVT == MVT::i64))
+static SDValue combineMulWide(SDNode *N, TargetLowering::DAGCombinerInfo &DCI,
+ CodeGenOptLevel OptLevel) {
+ if (OptLevel == CodeGenOptLevel::None)
return SDValue();
+
SDValue Op = N->getOperand(0);
- if (!(Op.getOpcode() == ISD::MUL || Op.getOpcode() == ISD::SHL))
- return SDValue();
- if (Op.getOpcode() == ISD::SHL && !isa<ConstantSDNode>(Op.getOperand(1)))
+ if (!Op.hasOneUse())
return SDValue();
+ EVT ToVT = N->getValueType(0);
EVT FromVT = Op.getValueType();
- if (!(FromVT == MVT::i16 || FromVT == MVT::i32))
- return SDValue();
- if (ExtOpcode == ISD::SIGN_EXTEND && !Op->getFlags().hasNoSignedWrap())
+ if (!((ToVT == MVT::i32 && FromVT == MVT::i16) ||
+ (ToVT == MVT::i64 && FromVT == MVT::i32)))
return SDValue();
- if (ExtOpcode == ISD::ZERO_EXTEND && !Op->getFlags().hasNoUnsignedWrap())
- return SDValue();
- if (ExtOpcode == ISD::ANY_EXTEND && !Op->getFlags().hasNoSignedWrap() &&
- !Op->getFlags().hasNoUnsignedWrap())
+ if (!(Op.getOpcode() == ISD::MUL ||
+ (Op.getOpcode() == ISD::SHL && isa<ConstantSDNode>(Op.getOperand(1)))))
return SDValue();
SDLoc DL(N);
+ unsigned ExtOpcode = N->getOpcode();
unsigned Opcode = 0;
- if (ExtOpcode == ISD::SIGN_EXTEND)
+ if (ExtOpcode == ISD::SIGN_EXTEND && Op->getFlags().hasNoSignedWrap())
Opcode = NVPTXISD::MUL_WIDE_SIGNED;
- else if (ExtOpcode == ISD::ZERO_EXTEND)
+ else if (ExtOpcode == ISD::ZERO_EXTEND && Op->getFlags().hasNoUnsignedWrap())
Opcode = NVPTXISD::MUL_WIDE_UNSIGNED;
- else if (ExtOpcode == ISD::ANY_EXTEND && Op->getFlags().hasNoUnsignedWrap())
- Opcode = NVPTXISD::MUL_WIDE_UNSIGNED;
- else if (ExtOpcode == ISD::ANY_EXTEND && Op->getFlags().hasNoSignedWrap())
- Opcode = NVPTXISD::MUL_WIDE_SIGNED;
else
- assert(false);
+ return SDValue();
SDValue RHS = Op.getOperand(1);
if (Op.getOpcode() == ISD::SHL) {
- const auto ShiftAmt =
- cast<ConstantSDNode>(Op.getOperand(1))->getZExtValue();
+ const auto ShiftAmt = Op.getConstantOperandVal(1);
const auto MulVal = APInt(ToVT.getSizeInBits(), 1) << ShiftAmt;
RHS = DCI.DAG.getConstant(MulVal, DL, ToVT);
}
@@ -5977,10 +5967,9 @@ SDValue NVPTXTargetLowering::PerformDAGCombine(SDNode *N,
return combineADDRSPACECAST(N, DCI);
case ISD::AND:
return PerformANDCombine(N, DCI);
- case ISD::ANY_EXTEND:
case ISD::SIGN_EXTEND:
case ISD::ZERO_EXTEND:
- return PerformExtendMULWIDECombine(N, DCI);
+ return combineMulWide(N, DCI, OptLevel);
case ISD::BUILD_VECTOR:
return PerformBUILD_VECTORCombine(N, DCI);
case ISD::EXTRACT_VECTOR_ELT:
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index b353ab26afc88..dbfe798b4f0a6 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -125,8 +125,6 @@ def doF32FTZ : Predicate<"useF32FTZ()">;
def doNoF32FTZ : Predicate<"!useF32FTZ()">;
def doRsqrtOpt : Predicate<"doRsqrtOpt()">;
-def doMulWide : Predicate<"doMulWide">;
-
def hasHWROT32 : Predicate<"Subtarget->hasHWROT32()">;
def noHWROT32 : Predicate<"!Subtarget->hasHWROT32()">;
def hasDotInstructions : Predicate<"Subtarget->hasDotInstructions()">;
@@ -860,11 +858,11 @@ def MULWIDEU32Imm32 :
BasicNVPTXInst<(outs B32:$dst), (ins B16:$a, i32imm:$b), "mul.wide.u16">;
def SDTMulWide : SDTypeProfile<1, 2, [SDTCisInt<0>, SDTCisInt<1>, SDTCisSameAs<1, 2>]>;
-def mul_wide_signed : SDNode<"NVPTXISD::MUL_WIDE_SIGNED", SDTMulWide>;
-def mul_wide_unsigned : SDNode<"NVPTXISD::MUL_WIDE_UNSIGNED", SDTMulWide>;
+def mul_wide_signed : SDNode<"NVPTXISD::MUL_WIDE_SIGNED", SDTMulWide, [SDNPCommutative]>;
+def mul_wide_unsigned : SDNode<"NVPTXISD::MUL_WIDE_UNSIGNED", SDTMulWide, [SDNPCommutative]>;
// Matchers for signed, unsigned mul.wide ISD nodes.
-let Predicates = [doMulWide] in {
+let Predicates = [hasOptEnabled] in {
def : Pat<(i32 (mul_wide_signed i16:$a, i16:$b)), (MULWIDES32 $a, $b)>;
def : Pat<(i32 (mul_wide_signed i16:$a, imm:$b)), (MULWIDES32Imm $a, imm:$b)>;
def : Pat<(i32 (mul_wide_unsigned i16:$a, i16:$b)), (MULWIDEU32 $a, $b)>;
@@ -876,85 +874,6 @@ let Predicates = [doMulWide] in {
def : Pat<(i64 (mul_wide_unsigned i32:$a, imm:$b)), (MULWIDEU64Imm $a, imm:$b)>;
}
-// Predicates used for converting some patterns to mul.wide.
-def SInt32Const : PatLeaf<(imm), [{
- const APInt &v = N->getAPIntValue();
- return v.isSignedIntN(32);
-}]>;
-
-def UInt32Const : PatLeaf<(imm), [{
- const APInt &v = N->getAPIntValue();
- return v.isIntN(32);
-}]>;
-
-def SInt16Const : PatLeaf<(imm), [{
- const APInt &v = N->getAPIntValue();
- return v.isSignedIntN(16);
-}]>;
-
-def UInt16Const : PatLeaf<(imm), [{
- const APInt &v = N->getAPIntValue();
- return v.isIntN(16);
-}]>;
-
-def IntConst_0_30 : PatLeaf<(imm), [{
- // Check if 0 <= v < 31; only then will the result of (x << v) be an int32.
- const APInt &v = N->getAPIntValue();
- return v.sge(0) && v.slt(31);
-}]>;
-
-def IntConst_0_14 : PatLeaf<(imm), [{
- // Check if 0 <= v < 15; only then will the result of (x << v) be an int16.
- const APInt &v = N->getAPIntValue();
- return v.sge(0) && v.slt(15);
-}]>;
-
-def SHL2MUL32 : SDNodeXForm<imm, [{
- const APInt &v = N->getAPIntValue();
- APInt temp(32, 1);
- return CurDAG->getTargetConstant(temp.shl(v), SDLoc(N), MVT::i32);
-}]>;
-
-def SHL2MUL16 : SDNodeXForm<imm, [{
- const APInt &v = N->getAPIntValue();
- APInt temp(16, 1);
- return CurDAG->getTargetConstant(temp.shl(v), SDLoc(N), MVT::i16);
-}]>;
-
-// Convert "sign/zero-extend, then shift left by an immediate" to mul.wide.
-let Predicates = [doMulWide] in {
- def : Pat<(shl (sext i32:$a), (i32 IntConst_0_30:$b)),
- (MULWIDES64Imm $a, (SHL2MUL32 $b))>;
- def : Pat<(shl (zext i32:$a), (i32 IntConst_0_30:$b)),
- (MULWIDEU64Imm $a, (SHL2MUL32 $b))>;
-
- def : Pat<(shl (sext i16:$a), (i16 IntConst_0_14:$b)),
- (MULWIDES32Imm $a, (SHL2MUL16 $b))>;
- def : Pat<(shl (zext i16:$a), (i16 IntConst_0_14:$b)),
- (MULWIDEU32Imm $a, (SHL2MUL16 $b))>;
-
- // Convert "sign/zero-extend then multiply" to mul.wide.
- def : Pat<(mul (sext i32:$a), (sext i32:$b)),
- (MULWIDES64 $a, $b)>;
- def : Pat<(mul (sext i32:$a), (i64 SInt32Const:$b)),
- (MULWIDES64Imm64 $a, (i64 SInt32Const:$b))>;
-
- def : Pat<(mul (zext i32:$a), (zext i32:$b)),
- (MULWIDEU64 $a, $b)>;
- def : Pat<(mul (zext i32:$a), (i64 UInt32Const:$b)),
- (MULWIDEU64Imm64 $a, (i64 UInt32Const:$b))>;
-
- def : Pat<(mul (sext i16:$a), (sext i16:$b)),
- (MULWIDES32 $a, $b)>;
- def : Pat<(mul (sext i16:$a), (i32 SInt16Const:$b)),
- (MULWIDES32Imm32 $a, (i32 SInt16Const:$b))>;
-
- def : Pat<(mul (zext i16:$a), (zext i16:$b)),
- (MULWIDEU32 $a, $b)>;
- def : Pat<(mul (zext i16:$a), (i32 UInt16Const:$b)),
- (MULWIDEU32Imm32 $a, (i32 UInt16Const:$b))>;
-}
-
//
// Integer multiply-add
//
@@ -990,33 +909,38 @@ defm MAD32 : MAD<"mad.lo.s32", i32, B32, i32imm>;
defm MAD64 : MAD<"mad.lo.s64", i64, B64, i64imm>;
}
-multiclass MAD_WIDE<string PtxSuffix, SDNode Op, ValueType BigVT, NVPTXRegClass BigReg, Operand BigImm, ValueType SmallVT, NVPTXRegClass SmallReg, Operand SmallImm> {
+multiclass MAD_WIDE<string PtxSuffix, OneUse2 Op, RegTyInfo BigT, RegTyInfo SmallT> {
def rrr:
- BasicNVPTXInst<(outs BigReg:$dst),
- (ins SmallReg:$a, SmallReg:$b, BigReg:$c),
+ BasicNVPTXInst<(outs BigT.RC:$dst),
+ (ins SmallT.RC:$a, SmallT.RC:$b, BigT.RC:$c),
"mad.wide." # PtxSuffix,
- [(set BigVT:$dst, (add (Op SmallVT:$a, SmallVT:$b), BigVT:$c))]>;
+ [(set BigT.Ty:$dst, (add (Op SmallT.Ty:$a, SmallT.Ty:$b), BigT.Ty:$c))]>;
def rri:
- BasicNVPTXInst<(outs BigReg:$dst),
- (ins SmallReg:$a, SmallReg:$b, BigImm:$c),
+ BasicNVPTXInst<(outs BigT.RC:$dst),
+ (ins SmallT.RC:$a, SmallT.RC:$b, BigT.Imm:$c),
"mad.wide." # PtxSuffix,
- [(set BigVT:$dst, (add (Op SmallVT:$a, SmallVT:$b), imm:$c))]>;
+ [(set BigT.Ty:$dst, (add (Op SmallT.Ty:$a, SmallT.Ty:$b), imm:$c))]>;
def rir:
- BasicNVPTXInst<(outs BigReg:$dst),
- (ins SmallReg:$a, SmallImm:$b, BigReg:$c),
+ BasicNVPTXInst<(outs BigT.RC:$dst),
+ (ins SmallT.RC:$a, SmallT.Imm:$b, BigT.RC:$c),
"mad.wide." # PtxSuffix,
- [(set BigVT:$dst, (add (Op SmallVT:$a, imm:$b), BigVT:$c))]>;
+ [(set BigT.Ty:$dst, (add (Op SmallT.Ty:$a, imm:$b), BigT.Ty:$c))]>;
def rii:
- BasicNVPTXInst<(outs BigReg:$dst),
- (ins SmallReg:$a, SmallImm:$b, BigImm:$c),
+ BasicNVPTXInst<(outs BigT.RC:$dst),
+ (ins SmallT.RC:$a, SmallT.Imm:$b, BigT.Imm:$c),
"mad.wide." # PtxSuffix,
- [(set BigVT:$dst, (add (Op SmallVT:$a, imm:$b), imm:$c))]>;
+ [(set BigT.Ty:$dst, (add (Op SmallT.Ty:$a, imm:$b), imm:$c))]>;
}
-defm MAD_WIDE_U16 : MAD_WIDE<"u16", mul_wide_unsigned, i32, B32, i32imm, i16, B16, i16imm>;
-defm MAD_WIDE_S16 : MAD_WIDE<"s16", mul_wide_signed, i32, B32, i32imm, i16, B16, i16imm>;
-defm MAD_WIDE_U32 : MAD_WIDE<"u32", mul_wide_unsigned, i64, B64, i64imm, i32, B32, i32imm>;
-defm MAD_WIDE_S32 : MAD_WIDE<"s32", mul_wide_signed, i64, B64, i64imm, i32, B32, i32imm>;
+def mul_wide_unsigned_oneuse : OneUse2<mul_wide_unsigned>;
+def mul_wide_signed_oneuse : OneUse2<mul_wide_signed>;
+
+let Predicates = [hasOptEnabled] in {
+defm MAD_WIDE_U16 : MAD_WIDE<"u16", mul_wide_unsigned_oneuse, I32RT, I16RT>;
+defm MAD_WIDE_S16 : MAD_WIDE<"s16", mul_wide_signed_oneuse, I32RT, I16RT>;
+defm MAD_WIDE_U32 : MAD_WIDE<"u32", mul_wide_unsigned_oneuse, I64RT, I32RT>;
+defm MAD_WIDE_S32 : MAD_WIDE<"s32", mul_wide_signed_oneuse, I64RT, I32RT>;
+}
foreach t = [I16RT, I32RT, I64RT] in {
def NEG_S # t.Size :
diff --git a/llvm/test/CodeGen/NVPTX/combine-wide.ll b/llvm/test/CodeGen/NVPTX/combine-wide.ll
index 218ac4ed567f1..ed4a2b6e419c3 100644
--- a/llvm/test/CodeGen/NVPTX/combine-wide.ll
+++ b/llvm/test/CodeGen/NVPTX/combine-wide.ll
@@ -1,21 +1,38 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
-; RUN: llc < %s | FileCheck %s
+; RUN: llc < %s -O1 | FileCheck %s --check-prefixes=CHECK,O1
+; RUN: llc < %s -O0 | FileCheck %s --check-prefixes=CHECK,O0
target triple = "nvptx64-nvidia-cuda"
define i64 @t1(i32 %a, i32 %b, i64 %c) {
-; CHECK-LABEL: t1(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .b64 %rd<3>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r1, [t1_param_0];
-; CHECK-NEXT: ld.param.b32 %r2, [t1_param_1];
-; CHECK-NEXT: ld.param.b64 %rd1, [t1_param_2];
-; CHECK-NEXT: mad.wide.s32 %rd2, %r1, %r2, %rd1;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
-; CHECK-NEXT: ret;
+;
+; O1-LABEL: t1(
+; O1: {
+; O1-NEXT: .reg .b32 %r<3>;
+; O1-NEXT: .reg .b64 %rd<3>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b32 %r1, [t1_param_0];
+; O1-NEXT: ld.param.b32 %r2, [t1_param_1];
+; O1-NEXT: ld.param.b64 %rd1, [t1_param_2];
+; O1-NEXT: mad.wide.s32 %rd2, %r1, %r2, %rd1;
+; O1-NEXT: st.param.b64 [func_retval0], %rd2;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t1(
+; O0: {
+; O0-NEXT: .reg .b32 %r<4>;
+; O0-NEXT: .reg .b64 %rd<4>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b64 %rd1, [t1_param_2];
+; O0-NEXT: ld.param.b32 %r2, [t1_param_1];
+; O0-NEXT: ld.param.b32 %r1, [t1_param_0];
+; O0-NEXT: mul.lo.s32 %r3, %r1, %r2;
+; O0-NEXT: cvt.s64.s32 %rd2, %r3;
+; O0-NEXT: add.s64 %rd3, %rd1, %rd2;
+; O0-NEXT: st.param.b64 [func_retval0], %rd3;
+; O0-NEXT: ret;
%mul = mul nsw i32 %a, %b
%sext = sext i32 %mul to i64
%add = add i64 %c, %sext
@@ -23,18 +40,34 @@ define i64 @t1(i32 %a, i32 %b, i64 %c) {
}
define i64 @t2(i32 %a, i32 %b, i64 %c) {
-; CHECK-LABEL: t2(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .b64 %rd<3>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r1, [t2_param_0];
-; CHECK-NEXT: ld.param.b32 %r2, [t2_param_1];
-; CHECK-NEXT: ld.param.b64 %rd1, [t2_param_2];
-; CHECK-NEXT: mad.wide.s32 %rd2, %r1, %r2, %rd1;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
-; CHECK-NEXT: ret;
+;
+; O1-LABEL: t2(
+; O1: {
+; O1-NEXT: .reg .b32 %r<3>;
+; O1-NEXT: .reg .b64 %rd<3>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b32 %r1, [t2_param_0];
+; O1-NEXT: ld.param.b32 %r2, [t2_param_1];
+; O1-NEXT: ld.param.b64 %rd1, [t2_param_2];
+; O1-NEXT: mad.wide.s32 %rd2, %r1, %r2, %rd1;
+; O1-NEXT: st.param.b64 [func_retval0], %rd2;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t2(
+; O0: {
+; O0-NEXT: .reg .b32 %r<4>;
+; O0-NEXT: .reg .b64 %rd<4>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b64 %rd1, [t2_param_2];
+; O0-NEXT: ld.param.b32 %r2, [t2_param_1];
+; O0-NEXT: ld.param.b32 %r1, [t2_param_0];
+; O0-NEXT: mul.lo.s32 %r3, %r1, %r2;
+; O0-NEXT: cvt.s64.s32 %rd2, %r3;
+; O0-NEXT: add.s64 %rd3, %rd2, %rd1;
+; O0-NEXT: st.param.b64 [func_retval0], %rd3;
+; O0-NEXT: ret;
%mul = mul nsw i32 %a, %b
%sext = sext i32 %mul to i64
%add = add i64 %sext, %c
@@ -42,17 +75,32 @@ define i64 @t2(i32 %a, i32 %b, i64 %c) {
}
define i64 @t3(i32 %a, i32 %b) {
-; CHECK-LABEL: t3(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r1, [t3_param_0];
-; CHECK-NEXT: ld.param.b32 %r2, [t3_param_1];
-; CHECK-NEXT: mad.wide.s32 %rd1, %r1, %r2, 1;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
-; CHECK-NEXT: ret;
+;
+; O1-LABEL: t3(
+; O1: {
+; O1-NEXT: .reg .b32 %r<3>;
+; O1-NEXT: .reg .b64 %rd<2>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b32 %r1, [t3_param_0];
+; O1-NEXT: ld.param.b32 %r2, [t3_param_1];
+; O1-NEXT: mad.wide.s32 %rd1, %r1, %r2, 1;
+; O1-NEXT: st.param.b64 [func_retval0], %rd1;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t3(
+; O0: {
+; O0-NEXT: .reg .b32 %r<4>;
+; O0-NEXT: .reg .b64 %rd<3>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r2, [t3_param_1];
+; O0-NEXT: ld.param.b32 %r1, [t3_param_0];
+; O0-NEXT: mul.lo.s32 %r3, %r1, %r2;
+; O0-NEXT: cvt.s64.s32 %rd1, %r3;
+; O0-NEXT: add.s64 %rd2, %rd1, 1;
+; O0-NEXT: st.param.b64 [func_retval0], %rd2;
+; O0-NEXT: ret;
%mul = mul nsw i32 %a, %b
%sext = sext i32 %mul to i64
%add = add i64 1, %sext
@@ -60,17 +108,32 @@ define i64 @t3(i32 %a, i32 %b) {
}
define i64 @t4(i32 %a, i64 %c) {
-; CHECK-LABEL: t4(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .b64 %rd<3>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r1, [t4_param_0];
-; CHECK-NEXT: ld.param.b64 %rd1, [t4_param_1];
-; CHECK-NEXT: mad.wide.s32 %rd2, %r1, 3, %rd1;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
-; CHECK-NEXT: ret;
+;
+; O1-LABEL: t4(
+; O1: {
+; O1-NEXT: .reg .b32 %r<2>;
+; O1-NEXT: .reg .b64 %rd<3>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b32 %r1, [t4_param_0];
+; O1-NEXT: ld.param.b64 %rd1, [t4_param_1];
+; O1-NEXT: mad.wide.s32 %rd2, %r1, 3, %rd1;
+; O1-NEXT: st.param.b64 [func_retval0], %rd2;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t4(
+; O0: {
+; O0-NEXT: .reg .b32 %r<3>;
+; O0-NEXT: .reg .b64 %rd<4>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b64 %rd1, [t4_param_1];
+; O0-NEXT: ld.param.b32 %r1, [t4_param_0];
+; O0-NEXT: mul.lo.s32 %r2, %r1, 3;
+; O0-NEXT: cvt.s64.s32 %rd2, %r2;
+; O0-NEXT: add.s64 %rd3, %rd1, %rd2;
+; O0-NEXT: st.param.b64 [func_retval0], %rd3;
+; O0-NEXT: ret;
%mul = mul nsw i32 %a, 3
%sext = sext i32 %mul to i64
%add = add i64 %c, %sext
@@ -78,16 +141,30 @@ define i64 @t4(i32 %a, i64 %c) {
}
define i64 @t4_1(i32 %a, i64 %c) {
-; CHECK-LABEL: t4_1(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r1, [t4_1_param_0];
-; CHECK-NEXT: mad.wide.s32 %rd1, %r1, 3, 5;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
-; CHECK-NEXT: ret;
+;
+; O1-LABEL: t4_1(
+; O1: {
+; O1-NEXT: .reg .b32 %r<2>;
+; O1-NEXT: .reg .b64 %rd<2>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b32 %r1, [t4_1_param_0];
+; O1-NEXT: mad.wide.s32 %rd1, %r1, 3, 5;
+; O1-NEXT: st.param.b64 [func_retval0], %rd1;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t4_1(
+; O0: {
+; O0-NEXT: .reg .b32 %r<3>;
+; O0-NEXT: .reg .b64 %rd<3>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r1, [t4_1_param_0];
+; O0-NEXT: mul.lo.s32 %r2, %r1, 3;
+; O0-NEXT: cvt.s64.s32 %rd1, %r2;
+; O0-NEXT: add.s64 %rd2, %rd1, 5;
+; O0-NEXT: st.param.b64 [func_retval0], %rd2;
+; O0-NEXT: ret;
%mul = mul nsw i32 %a, 3
%sext = sext i32 %mul to i64
%add = add i64 5, %sext
@@ -95,18 +172,34 @@ define i64 @t4_1(i32 %a, i64 %c) {
}
define i64 @t5(i32 %a, i32 %b, i64 %c) {
-; CHECK-LABEL: t5(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .b64 %rd<3>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r1, [t5_param_0];
-; CHECK-NEXT: ld.param.b32 %r2, [t5_param_1];
-; CHECK-NEXT: ld.param.b64 %rd1, [t5_param_2];
-; CHECK-NEXT: mad.wide.u32 %rd2, %r1, %r2, %rd1;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
-; CHECK-NEXT: ret;
+;
+; O1-LABEL: t5(
+; O1: {
+; O1-NEXT: .reg .b32 %r<3>;
+; O1-NEXT: .reg .b64 %rd<3>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b32 %r1, [t5_param_0];
+; O1-NEXT: ld.param.b32 %r2, [t5_param_1];
+; O1-NEXT: ld.param.b64 %rd1, [t5_param_2];
+; O1-NEXT: mad.wide.u32 %rd2, %r1, %r2, %rd1;
+; O1-NEXT: st.param.b64 [func_retval0], %rd2;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t5(
+; O0: {
+; O0-NEXT: .reg .b32 %r<4>;
+; O0-NEXT: .reg .b64 %rd<4>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b64 %rd1, [t5_param_2];
+; O0-NEXT: ld.param.b32 %r2, [t5_param_1];
+; O0-NEXT: ld.param.b32 %r1, [t5_param_0];
+; O0-NEXT: mul.lo.s32 %r3, %r1, %r2;
+; O0-NEXT: cvt.u64.u32 %rd2, %r3;
+; O0-NEXT: add.s64 %rd3, %rd1, %rd2;
+; O0-NEXT: st.param.b64 [func_retval0], %rd3;
+; O0-NEXT: ret;
%mul = mul nuw i32 %a, %b
%zext = zext i32 %mul to i64
%add = add i64 %c, %zext
@@ -114,18 +207,34 @@ define i64 @t5(i32 %a, i32 %b, i64 %c) {
}
define i64 @t6(i32 %a, i32 %b, i64 %c) {
-; CHECK-LABEL: t6(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .b64 %rd<3>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r1, [t6_param_0];
-; CHECK-NEXT: ld.param.b32 %r2, [t6_param_1];
-; CHECK-NEXT: ld.param.b64 %rd1, [t6_param_2];
-; CHECK-NEXT: mad.wide.u32 %rd2, %r1, %r2, %rd1;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
-; CHECK-NEXT: ret;
+;
+; O1-LABEL: t6(
+; O1: {
+; O1-NEXT: .reg .b32 %r<3>;
+; O1-NEXT: .reg .b64 %rd<3>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b32 %r1, [t6_param_0];
+; O1-NEXT: ld.param.b32 %r2, [t6_param_1];
+; O1-NEXT: ld.param.b64 %rd1, [t6_param_2];
+; O1-NEXT: mad.wide.u32 %rd2, %r1, %r2, %rd1;
+; O1-NEXT: st.param.b64 [func_retval0], %rd2;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t6(
+; O0: {
+; O0-NEXT: .reg .b32 %r<4>;
+; O0-NEXT: .reg .b64 %rd<4>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b64 %rd1, [t6_param_2];
+; O0-NEXT: ld.param.b32 %r2, [t6_param_1];
+; O0-NEXT: ld.param.b32 %r1, [t6_param_0];
+; O0-NEXT: mul.lo.s32 %r3, %r1, %r2;
+; O0-NEXT: cvt.u64.u32 %rd2, %r3;
+; O0-NEXT: add.s64 %rd3, %rd2, %rd1;
+; O0-NEXT: st.param.b64 [func_retval0], %rd3;
+; O0-NEXT: ret;
%mul = mul nuw i32 %a, %b
%zext = zext i32 %mul to i64
%add = add i64 %zext, %c
@@ -133,236 +242,419 @@ define i64 @t6(i32 %a, i32 %b, i64 %c) {
}
define i32 @t7(i16 %a, i16 %b) {
-; CHECK-LABEL: t7(
-; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<4>;
-; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b16 %rs1, [t7_param_0];
-; CHECK-NEXT: ld.param.b16 %rs2, [t7_param_1];
-; CHECK-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
-; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
-; CHECK-NEXT: ret;
+;
+; O1-LABEL: t7(
+; O1: {
+; O1-NEXT: .reg .b16 %rs<4>;
+; O1-NEXT: .reg .b32 %r<2>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b16 %rs1, [t7_param_0];
+; O1-NEXT: ld.param.b16 %rs2, [t7_param_1];
+; O1-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; O1-NEXT: cvt.u32.u16 %r1, %rs3;
+; O1-NEXT: st.param.b32 [func_retval0], %r1;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t7(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<4>;
+; O0-NEXT: .reg .b32 %r<2>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b16 %rs2, [t7_param_1];
+; O0-NEXT: ld.param.b16 %rs1, [t7_param_0];
+; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; O0-NEXT: cvt.u32.u16 %r1, %rs3;
+; O0-NEXT: st.param.b32 [func_retval0], %r1;
+; O0-NEXT: ret;
%mul = mul i16 %a, %b
%zext = zext i16 %mul to i32
ret i32 %zext
}
define i32 @t8(i16 %a, i16 %b) {
-; CHECK-LABEL: t8(
-; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<4>;
-; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b16 %rs1, [t8_param_0];
-; CHECK-NEXT: ld.param.b16 %rs2, [t8_param_1];
-; CHECK-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
-; CHECK-NEXT: cvt.s32.s16 %r1, %rs3;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
-; CHECK-NEXT: ret;
+;
+; O1-LABEL: t8(
+; O1: {
+; O1-NEXT: .reg .b16 %rs<4>;
+; O1-NEXT: .reg .b32 %r<2>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b16 %rs1, [t8_param_0];
+; O1-NEXT: ld.param.b16 %rs2, [t8_param_1];
+; O1-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; O1-NEXT: cvt.s32.s16 %r1, %rs3;
+; O1-NEXT: st.param.b32 [func_retval0], %r1;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t8(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<4>;
+; O0-NEXT: .reg .b32 %r<2>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b16 %rs2, [t8_param_1];
+; O0-NEXT: ld.param.b16 %rs1, [t8_param_0];
+; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; O0-NEXT: cvt.s32.s16 %r1, %rs3;
+; O0-NEXT: st.param.b32 [func_retval0], %r1;
+; O0-NEXT: ret;
%mul = mul i16 %a, %b
%sext = sext i16 %mul to i32
ret i32 %sext
}
define i64 @t9(i32 %a, i32 %b) {
-; CHECK-LABEL: t9(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<4>;
-; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r1, [t9_param_0];
-; CHECK-NEXT: ld.param.b32 %r2, [t9_param_1];
-; CHECK-NEXT: mul.lo.s32 %r3, %r1, %r2;
-; CHECK-NEXT: cvt.u64.u32 %rd1, %r3;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
-; CHECK-NEXT: ret;
+;
+; O1-LABEL: t9(
+; O1: {
+; O1-NEXT: .reg .b32 %r<4>;
+; O1-NEXT: .reg .b64 %rd<2>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b32 %r1, [t9_param_0];
+; O1-NEXT: ld.param.b32 %r2, [t9_param_1];
+; O1-NEXT: mul.lo.s32 %r3, %r1, %r2;
+; O1-NEXT: cvt.u64.u32 %rd1, %r3;
+; O1-NEXT: st.param.b64 [func_retval0], %rd1;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t9(
+; O0: {
+; O0-NEXT: .reg .b32 %r<4>;
+; O0-NEXT: .reg .b64 %rd<2>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r2, [t9_param_1];
+; O0-NEXT: ld.param.b32 %r1, [t9_param_0];
+; O0-NEXT: mul.lo.s32 %r3, %r1, %r2;
+; O0-NEXT: cvt.u64.u32 %rd1, %r3;
+; O0-NEXT: st.param.b64 [func_retval0], %rd1;
+; O0-NEXT: ret;
%mul = mul i32 %a, %b
%zext = zext i32 %mul to i64
ret i64 %zext
}
define i64 @t10(i32 %a, i32 %b) {
-; CHECK-LABEL: t10(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<4>;
-; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r1, [t10_param_0];
-; CHECK-NEXT: ld.param.b32 %r2, [t10_param_1];
-; CHECK-NEXT: mul.lo.s32 %r3, %r1, %r2;
-; CHECK-NEXT: cvt.s64.s32 %rd1, %r3;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
-; CHECK-NEXT: ret;
+;
+; O1-LABEL: t10(
+; O1: {
+; O1-NEXT: .reg .b32 %r<4>;
+; O1-NEXT: .reg .b64 %rd<2>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b32 %r1, [t10_param_0];
+; O1-NEXT: ld.param.b32 %r2, [t10_param_1];
+; O1-NEXT: mul.lo.s32 %r3, %r1, %r2;
+; O1-NEXT: cvt.s64.s32 %rd1, %r3;
+; O1-NEXT: st.param.b64 [func_retval0], %rd1;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t10(
+; O0: {
+; O0-NEXT: .reg .b32 %r<4>;
+; O0-NEXT: .reg .b64 %rd<2>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r2, [t10_param_1];
+; O0-NEXT: ld.param.b32 %r1, [t10_param_0];
+; O0-NEXT: mul.lo.s32 %r3, %r1, %r2;
+; O0-NEXT: cvt.s64.s32 %rd1, %r3;
+; O0-NEXT: st.param.b64 [func_retval0], %rd1;
+; O0-NEXT: ret;
%mul = mul i32 %a, %b
%sext = sext i32 %mul to i64
ret i64 %sext
}
define i32 @t11(i16 %a, i16 %b) {
-; CHECK-LABEL: t11(
-; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<4>;
-; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b16 %rs1, [t11_param_0];
-; CHECK-NEXT: ld.param.b16 %rs2, [t11_param_1];
-; CHECK-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
-; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
-; CHECK-NEXT: ret;
+;
+; O1-LABEL: t11(
+; O1: {
+; O1-NEXT: .reg .b16 %rs<4>;
+; O1-NEXT: .reg .b32 %r<2>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b16 %rs1, [t11_param_0];
+; O1-NEXT: ld.param.b16 %rs2, [t11_param_1];
+; O1-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; O1-NEXT: cvt.u32.u16 %r1, %rs3;
+; O1-NEXT: st.param.b32 [func_retval0], %r1;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t11(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<4>;
+; O0-NEXT: .reg .b32 %r<2>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b16 %rs2, [t11_param_1];
+; O0-NEXT: ld.param.b16 %rs1, [t11_param_0];
+; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; O0-NEXT: cvt.u32.u16 %r1, %rs3;
+; O0-NEXT: st.param.b32 [func_retval0], %r1;
+; O0-NEXT: ret;
%mul = mul nsw i16 %a, %b
%zext = zext i16 %mul to i32
ret i32 %zext
}
define i32 @t12(i16 %a, i16 %b) {
-; CHECK-LABEL: t12(
-; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<3>;
-; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b16 %rs1, [t12_param_0];
-; CHECK-NEXT: ld.param.b16 %rs2, [t12_param_1];
-; CHECK-NEXT: mul.wide.s16 %r1, %rs1, %rs2;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
-; CHECK-NEXT: ret;
+;
+; O1-LABEL: t12(
+; O1: {
+; O1-NEXT: .reg .b16 %rs<3>;
+; O1-NEXT: .reg .b32 %r<2>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b16 %rs1, [t12_param_0];
+; O1-NEXT: ld.param.b16 %rs2, [t12_param_1];
+; O1-NEXT: mul.wide.s16 %r1, %rs1, %rs2;
+; O1-NEXT: st.param.b32 [func_retval0], %r1;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t12(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<4>;
+; O0-NEXT: .reg .b32 %r<2>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b16 %rs2, [t12_param_1];
+; O0-NEXT: ld.param.b16 %rs1, [t12_param_0];
+; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; O0-NEXT: cvt.s32.s16 %r1, %rs3;
+; O0-NEXT: st.param.b32 [func_retval0], %r1;
+; O0-NEXT: ret;
%mul = mul nsw i16 %a, %b
%sext = sext i16 %mul to i32
ret i32 %sext
}
define i64 @t13(i32 %a, i32 %b) {
-; CHECK-LABEL: t13(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<4>;
-; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r1, [t13_param_0];
-; CHECK-NEXT: ld.param.b32 %r2, [t13_param_1];
-; CHECK-NEXT: mul.lo.s32 %r3, %r1, %r2;
-; CHECK-NEXT: cvt.u64.u32 %rd1, %r3;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
-; CHECK-NEXT: ret;
+;
+; O1-LABEL: t13(
+; O1: {
+; O1-NEXT: .reg .b32 %r<4>;
+; O1-NEXT: .reg .b64 %rd<2>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b32 %r1, [t13_param_0];
+; O1-NEXT: ld.param.b32 %r2, [t13_param_1];
+; O1-NEXT: mul.lo.s32 %r3, %r1, %r2;
+; O1-NEXT: cvt.u64.u32 %rd1, %r3;
+; O1-NEXT: st.param.b64 [func_retval0], %rd1;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t13(
+; O0: {
+; O0-NEXT: .reg .b32 %r<4>;
+; O0-NEXT: .reg .b64 %rd<2>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r2, [t13_param_1];
+; O0-NEXT: ld.param.b32 %r1, [t13_param_0];
+; O0-NEXT: mul.lo.s32 %r3, %r1, %r2;
+; O0-NEXT: cvt.u64.u32 %rd1, %r3;
+; O0-NEXT: st.param.b64 [func_retval0], %rd1;
+; O0-NEXT: ret;
%mul = mul nsw i32 %a, %b
%zext = zext i32 %mul to i64
ret i64 %zext
}
define i64 @t14(i32 %a, i32 %b) {
-; CHECK-LABEL: t14(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r1, [t14_param_0];
-; CHECK-NEXT: ld.param.b32 %r2, [t14_param_1];
-; CHECK-NEXT: mul.wide.s32 %rd1, %r1, %r2;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
-; CHECK-NEXT: ret;
+;
+; O1-LABEL: t14(
+; O1: {
+; O1-NEXT: .reg .b32 %r<3>;
+; O1-NEXT: .reg .b64 %rd<2>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b32 %r1, [t14_param_0];
+; O1-NEXT: ld.param.b32 %r2, [t14_param_1];
+; O1-NEXT: mul.wide.s32 %rd1, %r1, %r2;
+; O1-NEXT: st.param.b64 [func_retval0], %rd1;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t14(
+; O0: {
+; O0-NEXT: .reg .b32 %r<4>;
+; O0-NEXT: .reg .b64 %rd<2>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r2, [t14_param_1];
+; O0-NEXT: ld.param.b32 %r1, [t14_param_0];
+; O0-NEXT: mul.lo.s32 %r3, %r1, %r2;
+; O0-NEXT: cvt.s64.s32 %rd1, %r3;
+; O0-NEXT: st.param.b64 [func_retval0], %rd1;
+; O0-NEXT: ret;
%mul = mul nsw i32 %a, %b
%sext = sext i32 %mul to i64
ret i64 %sext
}
define i32 @t15(i16 %a, i16 %b) {
-; CHECK-LABEL: t15(
-; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<3>;
-; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b16 %rs1, [t15_param_0];
-; CHECK-NEXT: ld.param.b16 %rs2, [t15_param_1];
-; CHECK-NEXT: mul.wide.u16 %r1, %rs1, %rs2;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
-; CHECK-NEXT: ret;
+;
+; O1-LABEL: t15(
+; O1: {
+; O1-NEXT: .reg .b16 %rs<3>;
+; O1-NEXT: .reg .b32 %r<2>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b16 %rs1, [t15_param_0];
+; O1-NEXT: ld.param.b16 %rs2, [t15_param_1];
+; O1-NEXT: mul.wide.u16 %r1, %rs1, %rs2;
+; O1-NEXT: st.param.b32 [func_retval0], %r1;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t15(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<4>;
+; O0-NEXT: .reg .b32 %r<2>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b16 %rs2, [t15_param_1];
+; O0-NEXT: ld.param.b16 %rs1, [t15_param_0];
+; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; O0-NEXT: cvt.u32.u16 %r1, %rs3;
+; O0-NEXT: st.param.b32 [func_retval0], %r1;
+; O0-NEXT: ret;
%mul = mul nuw i16 %a, %b
%zext = zext i16 %mul to i32
ret i32 %zext
}
define i32 @t16(i16 %a, i16 %b) {
-; CHECK-LABEL: t16(
-; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<4>;
-; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b16 %rs1, [t16_param_0];
-; CHECK-NEXT: ld.param.b16 %rs2, [t16_param_1];
-; CHECK-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
-; CHECK-NEXT: cvt.s32.s16 %r1, %rs3;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
-; CHECK-NEXT: ret;
+;
+; O1-LABEL: t16(
+; O1: {
+; O1-NEXT: .reg .b16 %rs<4>;
+; O1-NEXT: .reg .b32 %r<2>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b16 %rs1, [t16_param_0];
+; O1-NEXT: ld.param.b16 %rs2, [t16_param_1];
+; O1-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; O1-NEXT: cvt.s32.s16 %r1, %rs3;
+; O1-NEXT: st.param.b32 [func_retval0], %r1;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t16(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<4>;
+; O0-NEXT: .reg .b32 %r<2>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b16 %rs2, [t16_param_1];
+; O0-NEXT: ld.param.b16 %rs1, [t16_param_0];
+; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; O0-NEXT: cvt.s32.s16 %r1, %rs3;
+; O0-NEXT: st.param.b32 [func_retval0], %r1;
+; O0-NEXT: ret;
%mul = mul nuw i16 %a, %b
%sext = sext i16 %mul to i32
ret i32 %sext
}
define i64 @t17(i32 %a, i32 %b) {
-; CHECK-LABEL: t17(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r1, [t17_param_0];
-; CHECK-NEXT: ld.param.b32 %r2, [t17_param_1];
-; CHECK-NEXT: mul.wide.u32 %rd1, %r1, %r2;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
-; CHECK-NEXT: ret;
+;
+; O1-LABEL: t17(
+; O1: {
+; O1-NEXT: .reg .b32 %r<3>;
+; O1-NEXT: .reg .b64 %rd<2>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b32 %r1, [t17_param_0];
+; O1-NEXT: ld.param.b32 %r2, [t17_param_1];
+; O1-NEXT: mul.wide.u32 %rd1, %r1, %r2;
+; O1-NEXT: st.param.b64 [func_retval0], %rd1;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t17(
+; O0: {
+; O0-NEXT: .reg .b32 %r<4>;
+; O0-NEXT: .reg .b64 %rd<2>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r2, [t17_param_1];
+; O0-NEXT: ld.param.b32 %r1, [t17_param_0];
+; O0-NEXT: mul.lo.s32 %r3, %r1, %r2;
+; O0-NEXT: cvt.u64.u32 %rd1, %r3;
+; O0-NEXT: st.param.b64 [func_retval0], %rd1;
+; O0-NEXT: ret;
%mul = mul nuw i32 %a, %b
%zext = zext i32 %mul to i64
ret i64 %zext
}
define i64 @t18(i32 %a, i32 %b) {
-; CHECK-LABEL: t18(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<4>;
-; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r1, [t18_param_0];
-; CHECK-NEXT: ld.param.b32 %r2, [t18_param_1];
-; CHECK-NEXT: mul.lo.s32 %r3, %r1, %r2;
-; CHECK-NEXT: cvt.s64.s32 %rd1, %r3;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
-; CHECK-NEXT: ret;
+;
+; O1-LABEL: t18(
+; O1: {
+; O1-NEXT: .reg .b32 %r<4>;
+; O1-NEXT: .reg .b64 %rd<2>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b32 %r1, [t18_param_0];
+; O1-NEXT: ld.param.b32 %r2, [t18_param_1];
+; O1-NEXT: mul.lo.s32 %r3, %r1, %r2;
+; O1-NEXT: cvt.s64.s32 %rd1, %r3;
+; O1-NEXT: st.param.b64 [func_retval0], %rd1;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t18(
+; O0: {
+; O0-NEXT: .reg .b32 %r<4>;
+; O0-NEXT: .reg .b64 %rd<2>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r2, [t18_param_1];
+; O0-NEXT: ld.param.b32 %r1, [t18_param_0];
+; O0-NEXT: mul.lo.s32 %r3, %r1, %r2;
+; O0-NEXT: cvt.s64.s32 %rd1, %r3;
+; O0-NEXT: st.param.b64 [func_retval0], %rd1;
+; O0-NEXT: ret;
%mul = mul nuw i32 %a, %b
%sext = sext i32 %mul to i64
ret i64 %sext
}
define i32 @t19(i16 %a, i16 %b) {
-; CHECK-LABEL: t19(
-; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<4>;
-; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b16 %rs1, [t19_param_0];
-; CHECK-NEXT: ld.param.b16 %rs2, [t19_param_1];
-; CHECK-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
-; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
-; CHECK-NEXT: ret;
+;
+; O1-LABEL: t19(
+; O1: {
+; O1-NEXT: .reg .b16 %rs<4>;
+; O1-NEXT: .reg .b32 %r<2>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b16 %rs1, [t19_param_0];
+; O1-NEXT: ld.param.b16 %rs2, [t19_param_1];
+; O1-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; O1-NEXT: cvt.u32.u16 %r1, %rs3;
+; O1-NEXT: st.param.b32 [func_retval0], %r1;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t19(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<4>;
+; O0-NEXT: .reg .b32 %r<2>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b16 %rs2, [t19_param_1];
+; O0-NEXT: ld.param.b16 %rs1, [t19_param_0];
+; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; O0-NEXT: cvt.u32.u16 %r1, %rs3;
+; O0-NEXT: st.param.b32 [func_retval0], %r1;
+; O0-NEXT: ret;
%mul = mul i16 %a, %b
%zext = zext i16 %mul to i32
ret i32 %zext
}
define i32 @t20(i16 %a) {
+;
; CHECK-LABEL: t20(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<3>;
@@ -380,6 +672,7 @@ define i32 @t20(i16 %a) {
}
define i64 @t21(i32 %a) {
+;
; CHECK-LABEL: t21(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<3>;
@@ -397,6 +690,7 @@ define i64 @t21(i32 %a) {
}
define i64 @t22(i32 %a) {
+;
; CHECK-LABEL: t22(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<3>;
@@ -414,6 +708,7 @@ define i64 @t22(i32 %a) {
}
define i32 @t23(i16 %a, i16 %b) {
+;
; CHECK-LABEL: t23(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<3>;
@@ -431,22 +726,36 @@ define i32 @t23(i16 %a, i16 %b) {
}
define i32 @t24(i16 %a, i16 %b) {
-; CHECK-LABEL: t24(
-; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<2>;
-; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b16 %rs1, [t24_param_0];
-; CHECK-NEXT: mul.wide.s16 %r1, %rs1, 16;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
-; CHECK-NEXT: ret;
+;
+; O1-LABEL: t24(
+; O1: {
+; O1-NEXT: .reg .b16 %rs<2>;
+; O1-NEXT: .reg .b32 %r<2>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b16 %rs1, [t24_param_0];
+; O1-NEXT: mul.wide.s16 %r1, %rs1, 16;
+; O1-NEXT: st.param.b32 [func_retval0], %r1;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t24(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<3>;
+; O0-NEXT: .reg .b32 %r<2>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b16 %rs1, [t24_param_0];
+; O0-NEXT: shl.b16 %rs2, %rs1, 4;
+; O0-NEXT: cvt.s32.s16 %r1, %rs2;
+; O0-NEXT: st.param.b32 [func_retval0], %r1;
+; O0-NEXT: ret;
%mul = shl nsw i16 %a, 4
%sext = sext i16 %mul to i32
ret i32 %sext
}
define i64 @t25(i32 %a) {
+;
; CHECK-LABEL: t25(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<3>;
@@ -464,38 +773,65 @@ define i64 @t25(i32 %a) {
}
define i64 @t26(i32 %a) {
-; CHECK-LABEL: t26(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r1, [t26_param_0];
-; CHECK-NEXT: mul.wide.s32 %rd1, %r1, 16;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
-; CHECK-NEXT: ret;
+;
+; O1-LABEL: t26(
+; O1: {
+; O1-NEXT: .reg .b32 %r<2>;
+; O1-NEXT: .reg .b64 %rd<2>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b32 %r1, [t26_param_0];
+; O1-NEXT: mul.wide.s32 %rd1, %r1, 16;
+; O1-NEXT: st.param.b64 [func_retval0], %rd1;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t26(
+; O0: {
+; O0-NEXT: .reg .b32 %r<3>;
+; O0-NEXT: .reg .b64 %rd<2>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r1, [t26_param_0];
+; O0-NEXT: shl.b32 %r2, %r1, 4;
+; O0-NEXT: cvt.s64.s32 %rd1, %r2;
+; O0-NEXT: st.param.b64 [func_retval0], %rd1;
+; O0-NEXT: ret;
%mul = shl nsw i32 %a, 4
%sext = sext i32 %mul to i64
ret i64 %sext
}
define i32 @t27(i16 %a, i16 %b) {
-; CHECK-LABEL: t27(
-; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<2>;
-; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b16 %rs1, [t27_param_0];
-; CHECK-NEXT: mul.wide.u16 %r1, %rs1, 16;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
-; CHECK-NEXT: ret;
+;
+; O1-LABEL: t27(
+; O1: {
+; O1-NEXT: .reg .b16 %rs<2>;
+; O1-NEXT: .reg .b32 %r<2>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b16 %rs1, [t27_param_0];
+; O1-NEXT: mul.wide.u16 %r1, %rs1, 16;
+; O1-NEXT: st.param.b32 [func_retval0], %r1;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t27(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<3>;
+; O0-NEXT: .reg .b32 %r<2>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b16 %rs1, [t27_param_0];
+; O0-NEXT: shl.b16 %rs2, %rs1, 4;
+; O0-NEXT: cvt.u32.u16 %r1, %rs2;
+; O0-NEXT: st.param.b32 [func_retval0], %r1;
+; O0-NEXT: ret;
%mul = shl nuw i16 %a, 4
%zext = zext i16 %mul to i32
ret i32 %zext
}
define i32 @t28(i16 %a, i16 %b) {
+;
; CHECK-LABEL: t28(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<3>;
@@ -513,22 +849,36 @@ define i32 @t28(i16 %a, i16 %b) {
}
define i64 @t29(i32 %a) {
-; CHECK-LABEL: t29(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r1, [t29_param_0];
-; CHECK-NEXT: mul.wide.u32 %rd1, %r1, 16;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
-; CHECK-NEXT: ret;
+;
+; O1-LABEL: t29(
+; O1: {
+; O1-NEXT: .reg .b32 %r<2>;
+; O1-NEXT: .reg .b64 %rd<2>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b32 %r1, [t29_param_0];
+; O1-NEXT: mul.wide.u32 %rd1, %r1, 16;
+; O1-NEXT: st.param.b64 [func_retval0], %rd1;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t29(
+; O0: {
+; O0-NEXT: .reg .b32 %r<3>;
+; O0-NEXT: .reg .b64 %rd<2>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r1, [t29_param_0];
+; O0-NEXT: shl.b32 %r2, %r1, 4;
+; O0-NEXT: cvt.u64.u32 %rd1, %r2;
+; O0-NEXT: st.param.b64 [func_retval0], %rd1;
+; O0-NEXT: ret;
%mul = shl nuw i32 %a, 4
%zext = zext i32 %mul to i64
ret i64 %zext
}
define i64 @t30(i32 %a) {
+;
; CHECK-LABEL: t30(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<3>;
@@ -546,36 +896,66 @@ define i64 @t30(i32 %a) {
}
define i64 @t31(i32 %a, i32 %b) {
-; CHECK-LABEL: t31(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<4>;
-; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r1, [t31_param_0];
-; CHECK-NEXT: ld.param.b32 %r2, [t31_param_1];
-; CHECK-NEXT: shl.b32 %r3, %r1, %r2;
-; CHECK-NEXT: cvt.s64.s32 %rd1, %r3;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
-; CHECK-NEXT: ret;
+;
+; O1-LABEL: t31(
+; O1: {
+; O1-NEXT: .reg .b32 %r<4>;
+; O1-NEXT: .reg .b64 %rd<2>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b32 %r1, [t31_param_0];
+; O1-NEXT: ld.param.b32 %r2, [t31_param_1];
+; O1-NEXT: shl.b32 %r3, %r1, %r2;
+; O1-NEXT: cvt.s64.s32 %rd1, %r3;
+; O1-NEXT: st.param.b64 [func_retval0], %rd1;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t31(
+; O0: {
+; O0-NEXT: .reg .b32 %r<4>;
+; O0-NEXT: .reg .b64 %rd<2>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r2, [t31_param_1];
+; O0-NEXT: ld.param.b32 %r1, [t31_param_0];
+; O0-NEXT: shl.b32 %r3, %r1, %r2;
+; O0-NEXT: cvt.s64.s32 %rd1, %r3;
+; O0-NEXT: st.param.b64 [func_retval0], %rd1;
+; O0-NEXT: ret;
%mul = shl nuw i32 %a, %b
%sext = sext i32 %mul to i64
ret i64 %sext
}
define i32 @t32(i16 %a, i16 %b, i32 %c) {
-; CHECK-LABEL: t32(
-; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<3>;
-; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b16 %rs1, [t32_param_0];
-; CHECK-NEXT: ld.param.b16 %rs2, [t32_param_1];
-; CHECK-NEXT: ld.param.b32 %r1, [t32_param_2];
-; CHECK-NEXT: mad.wide.s16 %r2, %rs1, %rs2, %r1;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
-; CHECK-NEXT: ret;
+;
+; O1-LABEL: t32(
+; O1: {
+; O1-NEXT: .reg .b16 %rs<3>;
+; O1-NEXT: .reg .b32 %r<3>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b16 %rs1, [t32_param_0];
+; O1-NEXT: ld.param.b16 %rs2, [t32_param_1];
+; O1-NEXT: ld.param.b32 %r1, [t32_param_2];
+; O1-NEXT: mad.wide.s16 %r2, %rs1, %rs2, %r1;
+; O1-NEXT: st.param.b32 [func_retval0], %r2;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t32(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<4>;
+; O0-NEXT: .reg .b32 %r<4>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r1, [t32_param_2];
+; O0-NEXT: ld.param.b16 %rs2, [t32_param_1];
+; O0-NEXT: ld.param.b16 %rs1, [t32_param_0];
+; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; O0-NEXT: cvt.s32.s16 %r2, %rs3;
+; O0-NEXT: add.s32 %r3, %r1, %r2;
+; O0-NEXT: st.param.b32 [func_retval0], %r3;
+; O0-NEXT: ret;
%mul = mul nsw i16 %a, %b
%sext = sext i16 %mul to i32
%add = add i32 %c, %sext
@@ -583,18 +963,34 @@ define i32 @t32(i16 %a, i16 %b, i32 %c) {
}
define i32 @t33(i16 %a, i16 %b, i32 %c) {
-; CHECK-LABEL: t33(
-; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<3>;
-; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b16 %rs1, [t33_param_0];
-; CHECK-NEXT: ld.param.b16 %rs2, [t33_param_1];
-; CHECK-NEXT: ld.param.b32 %r1, [t33_param_2];
-; CHECK-NEXT: mad.wide.s16 %r2, %rs1, %rs2, %r1;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
-; CHECK-NEXT: ret;
+;
+; O1-LABEL: t33(
+; O1: {
+; O1-NEXT: .reg .b16 %rs<3>;
+; O1-NEXT: .reg .b32 %r<3>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b16 %rs1, [t33_param_0];
+; O1-NEXT: ld.param.b16 %rs2, [t33_param_1];
+; O1-NEXT: ld.param.b32 %r1, [t33_param_2];
+; O1-NEXT: mad.wide.s16 %r2, %rs1, %rs2, %r1;
+; O1-NEXT: st.param.b32 [func_retval0], %r2;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t33(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<4>;
+; O0-NEXT: .reg .b32 %r<4>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r1, [t33_param_2];
+; O0-NEXT: ld.param.b16 %rs2, [t33_param_1];
+; O0-NEXT: ld.param.b16 %rs1, [t33_param_0];
+; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; O0-NEXT: cvt.s32.s16 %r2, %rs3;
+; O0-NEXT: add.s32 %r3, %r1, %r2;
+; O0-NEXT: st.param.b32 [func_retval0], %r3;
+; O0-NEXT: ret;
%mul = mul nsw i16 %a, %b
%sext = sext i16 %mul to i32
%add = add i32 %c, %sext
@@ -602,17 +998,32 @@ define i32 @t33(i16 %a, i16 %b, i32 %c) {
}
define i32 @t34(i16 %a, i16 %b) {
-; CHECK-LABEL: t34(
-; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<3>;
-; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b16 %rs1, [t34_param_0];
-; CHECK-NEXT: ld.param.b16 %rs2, [t34_param_1];
-; CHECK-NEXT: mad.wide.s16 %r1, %rs1, %rs2, 1;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
-; CHECK-NEXT: ret;
+;
+; O1-LABEL: t34(
+; O1: {
+; O1-NEXT: .reg .b16 %rs<3>;
+; O1-NEXT: .reg .b32 %r<2>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b16 %rs1, [t34_param_0];
+; O1-NEXT: ld.param.b16 %rs2, [t34_param_1];
+; O1-NEXT: mad.wide.s16 %r1, %rs1, %rs2, 1;
+; O1-NEXT: st.param.b32 [func_retval0], %r1;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t34(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<4>;
+; O0-NEXT: .reg .b32 %r<3>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b16 %rs2, [t34_param_1];
+; O0-NEXT: ld.param.b16 %rs1, [t34_param_0];
+; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; O0-NEXT: cvt.s32.s16 %r1, %rs3;
+; O0-NEXT: add.s32 %r2, %r1, 1;
+; O0-NEXT: st.param.b32 [func_retval0], %r2;
+; O0-NEXT: ret;
%mul = mul nsw i16 %a, %b
%sext = sext i16 %mul to i32
%add = add i32 1, %sext
@@ -620,17 +1031,32 @@ define i32 @t34(i16 %a, i16 %b) {
}
define i32 @t35(i16 %a, i32 %c) {
-; CHECK-LABEL: t35(
-; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<2>;
-; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b16 %rs1, [t35_param_0];
-; CHECK-NEXT: ld.param.b32 %r1, [t35_param_1];
-; CHECK-NEXT: mad.wide.s16 %r2, %rs1, 3, %r1;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
-; CHECK-NEXT: ret;
+;
+; O1-LABEL: t35(
+; O1: {
+; O1-NEXT: .reg .b16 %rs<2>;
+; O1-NEXT: .reg .b32 %r<3>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b16 %rs1, [t35_param_0];
+; O1-NEXT: ld.param.b32 %r1, [t35_param_1];
+; O1-NEXT: mad.wide.s16 %r2, %rs1, 3, %r1;
+; O1-NEXT: st.param.b32 [func_retval0], %r2;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t35(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<3>;
+; O0-NEXT: .reg .b32 %r<4>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r1, [t35_param_1];
+; O0-NEXT: ld.param.b16 %rs1, [t35_param_0];
+; O0-NEXT: mul.lo.s16 %rs2, %rs1, 3;
+; O0-NEXT: cvt.s32.s16 %r2, %rs2;
+; O0-NEXT: add.s32 %r3, %r1, %r2;
+; O0-NEXT: st.param.b32 [func_retval0], %r3;
+; O0-NEXT: ret;
%mul = mul nsw i16 %a, 3
%sext = sext i16 %mul to i32
%add = add i32 %c, %sext
@@ -638,16 +1064,30 @@ define i32 @t35(i16 %a, i32 %c) {
}
define i32 @t36(i16 %a, i32 %c) {
-; CHECK-LABEL: t36(
-; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<2>;
-; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b16 %rs1, [t36_param_0];
-; CHECK-NEXT: mad.wide.s16 %r1, %rs1, 3, 5;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
-; CHECK-NEXT: ret;
+;
+; O1-LABEL: t36(
+; O1: {
+; O1-NEXT: .reg .b16 %rs<2>;
+; O1-NEXT: .reg .b32 %r<2>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b16 %rs1, [t36_param_0];
+; O1-NEXT: mad.wide.s16 %r1, %rs1, 3, 5;
+; O1-NEXT: st.param.b32 [func_retval0], %r1;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t36(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<3>;
+; O0-NEXT: .reg .b32 %r<3>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b16 %rs1, [t36_param_0];
+; O0-NEXT: mul.lo.s16 %rs2, %rs1, 3;
+; O0-NEXT: cvt.s32.s16 %r1, %rs2;
+; O0-NEXT: add.s32 %r2, %r1, 5;
+; O0-NEXT: st.param.b32 [func_retval0], %r2;
+; O0-NEXT: ret;
%mul = mul nsw i16 %a, 3
%sext = sext i16 %mul to i32
%add = add i32 5, %sext
@@ -655,18 +1095,34 @@ define i32 @t36(i16 %a, i32 %c) {
}
define i32 @t37(i16 %a, i16 %b, i32 %c) {
-; CHECK-LABEL: t37(
-; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<3>;
-; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b16 %rs1, [t37_param_0];
-; CHECK-NEXT: ld.param.b16 %rs2, [t37_param_1];
-; CHECK-NEXT: ld.param.b32 %r1, [t37_param_2];
-; CHECK-NEXT: mad.wide.u16 %r2, %rs1, %rs2, %r1;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
-; CHECK-NEXT: ret;
+;
+; O1-LABEL: t37(
+; O1: {
+; O1-NEXT: .reg .b16 %rs<3>;
+; O1-NEXT: .reg .b32 %r<3>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b16 %rs1, [t37_param_0];
+; O1-NEXT: ld.param.b16 %rs2, [t37_param_1];
+; O1-NEXT: ld.param.b32 %r1, [t37_param_2];
+; O1-NEXT: mad.wide.u16 %r2, %rs1, %rs2, %r1;
+; O1-NEXT: st.param.b32 [func_retval0], %r2;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t37(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<4>;
+; O0-NEXT: .reg .b32 %r<4>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r1, [t37_param_2];
+; O0-NEXT: ld.param.b16 %rs2, [t37_param_1];
+; O0-NEXT: ld.param.b16 %rs1, [t37_param_0];
+; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; O0-NEXT: cvt.u32.u16 %r2, %rs3;
+; O0-NEXT: add.s32 %r3, %r1, %r2;
+; O0-NEXT: st.param.b32 [func_retval0], %r3;
+; O0-NEXT: ret;
%mul = mul nuw i16 %a, %b
%zext = zext i16 %mul to i32
%add = add i32 %c, %zext
@@ -674,19 +1130,209 @@ define i32 @t37(i16 %a, i16 %b, i32 %c) {
}
define i32 @t38(i16 %a, i16 %b, i32 %c) {
-; CHECK-LABEL: t38(
-; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<3>;
-; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b16 %rs1, [t38_param_0];
-; CHECK-NEXT: ld.param.b16 %rs2, [t38_param_1];
-; CHECK-NEXT: ld.param.b32 %r1, [t38_param_2];
-; CHECK-NEXT: mad.wide.u16 %r2, %rs1, %rs2, %r1;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
-; CHECK-NEXT: ret;
+;
+; O1-LABEL: t38(
+; O1: {
+; O1-NEXT: .reg .b16 %rs<3>;
+; O1-NEXT: .reg .b32 %r<3>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b16 %rs1, [t38_param_0];
+; O1-NEXT: ld.param.b16 %rs2, [t38_param_1];
+; O1-NEXT: ld.param.b32 %r1, [t38_param_2];
+; O1-NEXT: mad.wide.u16 %r2, %rs1, %rs2, %r1;
+; O1-NEXT: st.param.b32 [func_retval0], %r2;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t38(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<4>;
+; O0-NEXT: .reg .b32 %r<4>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r1, [t38_param_2];
+; O0-NEXT: ld.param.b16 %rs2, [t38_param_1];
+; O0-NEXT: ld.param.b16 %rs1, [t38_param_0];
+; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; O0-NEXT: cvt.u32.u16 %r2, %rs3;
+; O0-NEXT: add.s32 %r3, %r2, %r1;
+; O0-NEXT: st.param.b32 [func_retval0], %r3;
+; O0-NEXT: ret;
+ %mul = mul nuw i16 %a, %b
+ %zext = zext i16 %mul to i32
+ %add = add i32 %zext, %c
+ ret i32 %add
+}
+
+define i64 @t39(i16 %a, i16 %b) {
+; O1-LABEL: t39(
+; O1: {
+; O1-NEXT: .reg .b16 %rs<4>;
+; O1-NEXT: .reg .b64 %rd<2>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b16 %rs1, [t39_param_0];
+; O1-NEXT: ld.param.b16 %rs2, [t39_param_1];
+; O1-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; O1-NEXT: cvt.u64.u16 %rd1, %rs3;
+; O1-NEXT: st.param.b64 [func_retval0], %rd1;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t39(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<4>;
+; O0-NEXT: .reg .b64 %rd<2>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b16 %rs2, [t39_param_1];
+; O0-NEXT: ld.param.b16 %rs1, [t39_param_0];
+; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; O0-NEXT: cvt.u64.u16 %rd1, %rs3;
+; O0-NEXT: st.param.b64 [func_retval0], %rd1;
+; O0-NEXT: ret;
+ %mul = mul i16 %a, %b
+ %zext = zext i16 %mul to i64
+ ret i64 %zext
+}
+
+define i64 @t40(i16 %a, i16 %b) {
+; O1-LABEL: t40(
+; O1: {
+; O1-NEXT: .reg .b16 %rs<4>;
+; O1-NEXT: .reg .b64 %rd<2>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b16 %rs1, [t40_param_0];
+; O1-NEXT: ld.param.b16 %rs2, [t40_param_1];
+; O1-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; O1-NEXT: cvt.u64.u16 %rd1, %rs3;
+; O1-NEXT: st.param.b64 [func_retval0], %rd1;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t40(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<4>;
+; O0-NEXT: .reg .b64 %rd<2>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b16 %rs2, [t40_param_1];
+; O0-NEXT: ld.param.b16 %rs1, [t40_param_0];
+; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; O0-NEXT: cvt.u64.u16 %rd1, %rs3;
+; O0-NEXT: st.param.b64 [func_retval0], %rd1;
+; O0-NEXT: ret;
+ %mul = mul nuw i16 %a, %b
+ %zext = zext i16 %mul to i64
+ ret i64 %zext
+}
+
+define i64 @t41(i16 %a, i16 %b) {
+; O1-LABEL: t41(
+; O1: {
+; O1-NEXT: .reg .b16 %rs<4>;
+; O1-NEXT: .reg .b64 %rd<2>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b16 %rs1, [t41_param_0];
+; O1-NEXT: ld.param.b16 %rs2, [t41_param_1];
+; O1-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; O1-NEXT: cvt.s64.s16 %rd1, %rs3;
+; O1-NEXT: st.param.b64 [func_retval0], %rd1;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t41(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<4>;
+; O0-NEXT: .reg .b64 %rd<2>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b16 %rs2, [t41_param_1];
+; O0-NEXT: ld.param.b16 %rs1, [t41_param_0];
+; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; O0-NEXT: cvt.s64.s16 %rd1, %rs3;
+; O0-NEXT: st.param.b64 [func_retval0], %rd1;
+; O0-NEXT: ret;
+ %mul = mul nsw i16 %a, %b
+ %sext = sext i16 %mul to i64
+ ret i64 %sext
+}
+
+define i32 @t42(i16 %a, i16 %b, ptr %ptr) {
+; O1-LABEL: t42(
+; O1: {
+; O1-NEXT: .reg .b16 %rs<4>;
+; O1-NEXT: .reg .b32 %r<2>;
+; O1-NEXT: .reg .b64 %rd<2>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b16 %rs1, [t42_param_0];
+; O1-NEXT: ld.param.b16 %rs2, [t42_param_1];
+; O1-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; O1-NEXT: ld.param.b64 %rd1, [t42_param_2];
+; O1-NEXT: st.b16 [%rd1], %rs3;
+; O1-NEXT: cvt.u32.u16 %r1, %rs3;
+; O1-NEXT: st.param.b32 [func_retval0], %r1;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t42(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<4>;
+; O0-NEXT: .reg .b32 %r<2>;
+; O0-NEXT: .reg .b64 %rd<2>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b64 %rd1, [t42_param_2];
+; O0-NEXT: ld.param.b16 %rs2, [t42_param_1];
+; O0-NEXT: ld.param.b16 %rs1, [t42_param_0];
+; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; O0-NEXT: st.b16 [%rd1], %rs3;
+; O0-NEXT: cvt.u32.u16 %r1, %rs3;
+; O0-NEXT: st.param.b32 [func_retval0], %r1;
+; O0-NEXT: ret;
+ %mul = mul nuw i16 %a, %b
+ store i16 %mul, ptr %ptr
+ %zext = zext i16 %mul to i32
+ ret i32 %zext
+}
+
+define i32 @t43(i16 %a, i16 %b, i32 %c, ptr %ptr) {
+; O1-LABEL: t43(
+; O1: {
+; O1-NEXT: .reg .b16 %rs<4>;
+; O1-NEXT: .reg .b32 %r<4>;
+; O1-NEXT: .reg .b64 %rd<2>;
+; O1-EMPTY:
+; O1-NEXT: // %bb.0:
+; O1-NEXT: ld.param.b16 %rs1, [t43_param_0];
+; O1-NEXT: ld.param.b16 %rs2, [t43_param_1];
+; O1-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; O1-NEXT: ld.param.b64 %rd1, [t43_param_3];
+; O1-NEXT: st.b16 [%rd1], %rs3;
+; O1-NEXT: ld.param.b32 %r1, [t43_param_2];
+; O1-NEXT: cvt.u32.u16 %r2, %rs3;
+; O1-NEXT: add.s32 %r3, %r2, %r1;
+; O1-NEXT: st.param.b32 [func_retval0], %r3;
+; O1-NEXT: ret;
+;
+; O0-LABEL: t43(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<4>;
+; O0-NEXT: .reg .b32 %r<4>;
+; O0-NEXT: .reg .b64 %rd<2>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b64 %rd1, [t43_param_3];
+; O0-NEXT: ld.param.b32 %r1, [t43_param_2];
+; O0-NEXT: ld.param.b16 %rs2, [t43_param_1];
+; O0-NEXT: ld.param.b16 %rs1, [t43_param_0];
+; O0-NEXT: mul.lo.s16 %rs3, %rs1, %rs2;
+; O0-NEXT: st.b16 [%rd1], %rs3;
+; O0-NEXT: cvt.u32.u16 %r2, %rs3;
+; O0-NEXT: add.s32 %r3, %r2, %r1;
+; O0-NEXT: st.param.b32 [func_retval0], %r3;
+; O0-NEXT: ret;
%mul = mul nuw i16 %a, %b
+ store i16 %mul, ptr %ptr
%zext = zext i16 %mul to i32
%add = add i32 %zext, %c
ret i32 %add
>From e551e13c44f22eeb275d1dd618e87f9880a22e48 Mon Sep 17 00:00:00 2001
From: Justin Fargnoli <jfargnoli at nvidia.com>
Date: Mon, 28 Jul 2025 23:08:13 +0000
Subject: [PATCH 6/6] Fixup setTargetDAGCombine
---
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 3 +--
1 file changed, 1 insertion(+), 2 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index b846e7258d10b..d0ee1f1d51db1 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -841,8 +841,7 @@ NVPTXTargetLowering::NVPTXTargetLowering(const NVPTXTargetMachine &TM,
setTargetDAGCombine({ISD::ADD, ISD::AND, ISD::EXTRACT_VECTOR_ELT, ISD::FADD,
ISD::MUL, ISD::SHL, ISD::SREM, ISD::UREM, ISD::VSELECT,
ISD::BUILD_VECTOR, ISD::ADDRSPACECAST, ISD::LOAD,
- ISD::STORE, ISD::ZERO_EXTEND, ISD::SIGN_EXTEND,
- ISD::ANY_EXTEND});
+ ISD::STORE, ISD::ZERO_EXTEND, ISD::SIGN_EXTEND});
// setcc for f16x2 and bf16x2 needs special handling to prevent
// legalizer's attempt to scalarize it due to v2i1 not being legal.
More information about the llvm-commits
mailing list