[clang] [llvm] [NVPTX] Support inline asm with 128-bit operand in NVPTX backend (PR #97113)
via cfe-commits
cfe-commits at lists.llvm.org
Fri Jun 28 15:05:36 PDT 2024
https://github.com/Chengjunp updated https://github.com/llvm/llvm-project/pull/97113
>From b9d056d5aba5185dcfb44ee1062bc6fbc36d002d Mon Sep 17 00:00:00 2001
From: chengjunp <chengjunp at nvidia.com>
Date: Tue, 25 Jun 2024 17:20:27 +0000
Subject: [PATCH 1/6] [NVPTX] Support inline asm with 128-bit operand in NVPTX
backend
---
clang/lib/Basic/Targets/NVPTX.h | 1 +
.../NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp | 3 +
llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp | 2 +
llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp | 68 +++++++
llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.h | 3 +-
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 94 +++++++++
llvm/lib/Target/NVPTX/NVPTXISelLowering.h | 8 +
llvm/lib/Target/NVPTX/NVPTXInstrInfo.cpp | 2 +
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td | 12 +-
llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp | 4 +
llvm/lib/Target/NVPTX/NVPTXRegisterInfo.td | 3 +
.../CodeGen/NVPTX/inline-asm-b128-test1.ll | 92 +++++++++
.../CodeGen/NVPTX/inline-asm-b128-test2.ll | 57 ++++++
.../CodeGen/NVPTX/inline-asm-b128-test3.ll | 179 ++++++++++++++++++
14 files changed, 525 insertions(+), 3 deletions(-)
create mode 100644 llvm/test/CodeGen/NVPTX/inline-asm-b128-test1.ll
create mode 100644 llvm/test/CodeGen/NVPTX/inline-asm-b128-test2.ll
create mode 100644 llvm/test/CodeGen/NVPTX/inline-asm-b128-test3.ll
diff --git a/clang/lib/Basic/Targets/NVPTX.h b/clang/lib/Basic/Targets/NVPTX.h
index f476d49047c01..7e9b6b34df636 100644
--- a/clang/lib/Basic/Targets/NVPTX.h
+++ b/clang/lib/Basic/Targets/NVPTX.h
@@ -105,6 +105,7 @@ class LLVM_LIBRARY_VISIBILITY NVPTXTargetInfo : public TargetInfo {
case 'l':
case 'f':
case 'd':
+ case 'q':
Info.setAllowsRegister();
return true;
}
diff --git a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp
index b7a20c351f5ff..380d878c1f532 100644
--- a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp
+++ b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp
@@ -60,6 +60,9 @@ void NVPTXInstPrinter::printRegName(raw_ostream &OS, MCRegister Reg) const {
case 6:
OS << "%fd";
break;
+ case 7:
+ OS << "%rq";
+ break;
}
unsigned VReg = Reg.id() & 0x0FFFFFFF;
diff --git a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
index ca077d41d36ba..1645261d74d06 100644
--- a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
@@ -315,6 +315,8 @@ unsigned NVPTXAsmPrinter::encodeVirtualRegister(unsigned Reg) {
Ret = (5 << 28);
} else if (RC == &NVPTX::Float64RegsRegClass) {
Ret = (6 << 28);
+ } else if (RC == &NVPTX::Int128RegsRegClass) {
+ Ret = (7 << 28);
} else {
report_fatal_error("Bad register class");
}
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
index 1e1cbb15e33d4..05706e200bda6 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
@@ -519,6 +519,20 @@ void NVPTXDAGToDAGISel::Select(SDNode *N) {
if (tryConstantFP(N))
return;
break;
+ case ISD::CopyToReg: {
+ if (N->getOperand(1).getValueType() == MVT::i128) {
+ SelectV2I64toI128(N);
+ return;
+ }
+ break;
+ }
+ case ISD::CopyFromReg: {
+ if(N->getOperand(1).getValueType() == MVT::i128){
+ SelectI128toV2I64(N);
+ return;
+ }
+ break;
+ }
default:
break;
}
@@ -3798,6 +3812,60 @@ bool NVPTXDAGToDAGISel::SelectInlineAsmMemoryOperand(
return true;
}
+void NVPTXDAGToDAGISel::SelectV2I64toI128(SDNode *N) {
+ // Lower a CopyToReg with two 64-bit inputs
+ // Dst:i128, lo:i64, hi:i64
+ //
+ // CopyToReg Dst, lo, hi;
+ //
+ // ==>
+ //
+ // tmp = V2I64toI128 {lo, hi};
+ // CopyToReg Dst, tmp;
+ SDValue Dst = N->getOperand(1);
+ SDValue Lo = N->getOperand(2);
+ SDValue Hi = N->getOperand(3);
+
+ SDLoc DL(N);
+ SDNode *Mov =
+ CurDAG->getMachineNode(NVPTX::V2I64toI128, DL, MVT::i128, {Lo, Hi});
+
+ SmallVector<EVT, 8> ResultsType(N->value_begin(), N->value_end());
+ SmallVector<SDValue, 8> NewOps(N->getNumOperands() - 1);
+ NewOps[0] = N->getOperand(0);
+ NewOps[1] = Dst;
+ NewOps[2] = SDValue(Mov, 0);
+ if (N->getNumOperands() == 5)
+ NewOps[3] = N->getOperand(4);
+ SDValue NewValue = CurDAG->getNode(ISD::CopyToReg, DL, ResultsType, NewOps);
+
+ ReplaceNode(N, NewValue.getNode());
+}
+
+void NVPTXDAGToDAGISel::SelectI128toV2I64(SDNode *N) {
+ // Lower CopyFromReg from a 128-bit regs to two 64-bit regs
+ // Dst:i128, Src:i128
+ //
+ // {lo, hi} = CopyFromReg Src
+ //
+ // ==>
+ //
+ // {lo, hi} = I128toV2I64 Src
+ //
+ SDValue Ch = N->getOperand(0);
+ SDValue Src = N->getOperand(1);
+ SDValue Glue = N->getOperand(2);
+ SDLoc DL(N);
+
+ // Add Glue and Ch to the operands and results to avoid break the execution order
+ SDNode *Mov = CurDAG->getMachineNode(
+ NVPTX::I128toV2I64, DL,
+ {MVT::i64, MVT::i64, Ch.getValueType(), Glue.getValueType()},
+ {Src, Ch, Glue});
+
+ ReplaceNode(N, Mov);
+}
+
/// GetConvertOpcode - Returns the CVT_ instruction opcode that implements a
/// conversion from \p SrcTy to \p DestTy.
unsigned NVPTXDAGToDAGISel::GetConvertOpcode(MVT DestTy, MVT SrcTy,
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.h b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.h
index c5524351f2ff9..49626d4051485 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.h
+++ b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.h
@@ -74,7 +74,8 @@ class LLVM_LIBRARY_VISIBILITY NVPTXDAGToDAGISel : public SelectionDAGISel {
bool SelectSETP_F16X2(SDNode *N);
bool SelectSETP_BF16X2(SDNode *N);
bool tryEXTRACT_VECTOR_ELEMENT(SDNode *N);
-
+ void SelectV2I64toI128(SDNode *N);
+ void SelectI128toV2I64(SDNode *N);
inline SDValue getI32Imm(unsigned Imm, const SDLoc &DL) {
return CurDAG->getTargetConstant(Imm, DL, MVT::i32);
}
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 76633a437fe71..927c3ef55dcb1 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -859,6 +859,10 @@ NVPTXTargetLowering::NVPTXTargetLowering(const NVPTXTargetMachine &TM,
setBF16OperationAction(Op, MVT::v2bf16, Legal, Expand);
}
+ // Custom lowering for inline asm with 128-bit operands
+ setOperationAction(ISD::CopyToReg, MVT::i128, Custom);
+ setOperationAction(ISD::CopyFromReg, MVT::i128, Custom);
+
// No FEXP2, FLOG2. The PTX ex2 and log2 functions are always approximate.
// No FPOW or FREM in PTX.
@@ -2804,6 +2808,8 @@ NVPTXTargetLowering::LowerOperation(SDValue Op, SelectionDAG &DAG) const {
return LowerVectorArith(Op, DAG);
case ISD::DYNAMIC_STACKALLOC:
return LowerDYNAMIC_STACKALLOC(Op, DAG);
+ case ISD::CopyToReg:
+ return LowerCopyToReg_128(Op, DAG);
default:
llvm_unreachable("Custom lowering not defined for operation");
}
@@ -3094,6 +3100,53 @@ SDValue NVPTXTargetLowering::LowerSTOREi1(SDValue Op, SelectionDAG &DAG) const {
return Result;
}
+SDValue NVPTXTargetLowering::LowerCopyToReg_128(SDValue Op,
+ SelectionDAG &DAG) const {
+ // Change the CopyToReg to take in two 64-bit operands instead of a 128-bit
+ // operand so that it can pass the legalization.
+
+ assert(Op.getOperand(1).getValueType() == MVT::i128 &&
+ "Custom lowering for 128-bit CopyToReg only");
+
+ SDNode *Node = Op.getNode();
+ SDLoc DL(Node);
+
+ SDValue Cast = DAG.getNode(ISD::BITCAST, DL, MVT::v2i64, Op->getOperand(2));
+ SDValue Lo = DAG.getNode(ISD::EXTRACT_VECTOR_ELT, DL, MVT::i64, Cast,
+ DAG.getIntPtrConstant(0, DL));
+ SDValue Hi = DAG.getNode(ISD::EXTRACT_VECTOR_ELT, DL, MVT::i64, Cast,
+ DAG.getIntPtrConstant(1, DL));
+
+ SmallVector<SDValue, 8> NewOps(Op->getNumOperands() + 1);
+ SmallVector<EVT, 8> ResultsType(Node->value_begin(), Node->value_end());
+
+ NewOps[0] = Op->getOperand(0); // Chain
+ NewOps[1] = Op->getOperand(1); // Dst Reg
+ NewOps[2] = Lo; // Lower 64-bit
+ NewOps[3] = Hi; // Higher 64-bit
+ if (Op.getNumOperands() == 4)
+ NewOps[4] = Op->getOperand(3); // Glue if exists
+
+ return DAG.getNode(ISD::CopyToReg, DL, ResultsType, NewOps);
+}
+
+unsigned NVPTXTargetLowering::getNumRegisters(LLVMContext &Context, EVT VT,
+ std::optional<MVT> RegisterVT = std::nullopt) const {
+ if(VT == MVT::i128 && RegisterVT == MVT::i128)
+ return 1;
+ return TargetLoweringBase::getNumRegisters(Context, VT, RegisterVT);
+}
+
+bool NVPTXTargetLowering::splitValueIntoRegisterParts(
+ SelectionDAG &DAG, const SDLoc &DL, SDValue Val, SDValue *Parts,
+ unsigned NumParts, MVT PartVT, std::optional<CallingConv::ID> CC) const {
+ if (Val.getValueType() == MVT::i128 && NumParts == 1) {
+ Parts[0] = Val;
+ return true;
+ }
+ return false;
+}
+
// This creates target external symbol for a function parameter.
// Name of the symbol is composed from its index and the function name.
// Negative index corresponds to special parameter (unsized array) used for
@@ -5152,6 +5205,7 @@ NVPTXTargetLowering::getConstraintType(StringRef Constraint) const {
case 'l':
case 'f':
case 'd':
+ case 'q':
case '0':
case 'N':
return C_RegisterClass;
@@ -5177,6 +5231,12 @@ NVPTXTargetLowering::getRegForInlineAsmConstraint(const TargetRegisterInfo *TRI,
case 'l':
case 'N':
return std::make_pair(0U, &NVPTX::Int64RegsRegClass);
+ case 'q': {
+ if (STI.getSmVersion() < 70)
+ report_fatal_error("Inline asm with 128 bit operands is only "
+ "supported for sm_70 and higher!");
+ return std::make_pair(0U, &NVPTX::Int128RegsRegClass);
+ }
case 'f':
return std::make_pair(0U, &NVPTX::Float32RegsRegClass);
case 'd':
@@ -6244,6 +6304,37 @@ static void ReplaceINTRINSIC_W_CHAIN(SDNode *N, SelectionDAG &DAG,
}
}
+static void ReplaceCopyFromReg_128(SDNode *N, SelectionDAG &DAG,
+ SmallVectorImpl<SDValue> &Results) {
+ // Change the CopyFromReg to output 2 64-bit results instead of a 128-bit result
+ // so that it can pass the legalization
+ SDLoc DL(N);
+ SDValue Chain = N->getOperand(0);
+ SDValue Reg = N->getOperand(1);
+ SDValue Glue = N->getOperand(2);
+
+ assert(Reg.getValueType() == MVT::i128 &&
+ "Custom lowering for CopyFromReg with 128-bit reg only");
+ SmallVector<EVT, 8> ResultsType(4);
+ SmallVector<SDValue, 8> NewOps(3);
+ ResultsType[0] = MVT::i64;
+ ResultsType[1] = MVT::i64;
+ ResultsType[2] = N->getValueType(1);
+ ResultsType[3] = N->getValueType(2);
+
+ NewOps[0] = Chain;
+ NewOps[1] = Reg;
+ NewOps[2] = Glue;
+
+ SDValue NewValue = DAG.getNode(ISD::CopyFromReg, DL, ResultsType, NewOps);
+ SDValue Pair = DAG.getNode(ISD::BUILD_PAIR, DL, MVT::i128,
+ {NewValue.getValue(0), NewValue.getValue(1)});
+
+ Results.push_back(Pair);
+ Results.push_back(NewValue.getValue(2));
+ Results.push_back(NewValue.getValue(3));
+}
+
void NVPTXTargetLowering::ReplaceNodeResults(
SDNode *N, SmallVectorImpl<SDValue> &Results, SelectionDAG &DAG) const {
switch (N->getOpcode()) {
@@ -6255,6 +6346,9 @@ void NVPTXTargetLowering::ReplaceNodeResults(
case ISD::INTRINSIC_W_CHAIN:
ReplaceINTRINSIC_W_CHAIN(N, DAG, Results);
return;
+ case ISD::CopyFromReg:
+ ReplaceCopyFromReg_128(N, DAG, Results);
+ return;
}
}
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
index e211286fcc556..63262961b363e 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
@@ -640,6 +640,14 @@ class NVPTXTargetLowering : public TargetLowering {
SDValue LowerVAARG(SDValue Op, SelectionDAG &DAG) const;
SDValue LowerVASTART(SDValue Op, SelectionDAG &DAG) const;
+ SDValue LowerCopyToReg_128(SDValue Op, SelectionDAG &DAG) const;
+ unsigned getNumRegisters(LLVMContext &Context, EVT VT,
+ std::optional<MVT> RegisterVT) const override;
+ bool
+ splitValueIntoRegisterParts(SelectionDAG &DAG, const SDLoc &DL, SDValue Val,
+ SDValue *Parts, unsigned NumParts, MVT PartVT,
+ std::optional<CallingConv::ID> CC) const override;
+
void ReplaceNodeResults(SDNode *N, SmallVectorImpl<SDValue> &Results,
SelectionDAG &DAG) const override;
SDValue PerformDAGCombine(SDNode *N, DAGCombinerInfo &DCI) const override;
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.cpp b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.cpp
index b0d792b5ee3fe..673858f92e7ce 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.cpp
@@ -51,6 +51,8 @@ void NVPTXInstrInfo::copyPhysReg(MachineBasicBlock &MBB,
} else if (DestRC == &NVPTX::Int64RegsRegClass) {
Op = (SrcRC == &NVPTX::Int64RegsRegClass ? NVPTX::IMOV64rr
: NVPTX::BITCONVERT_64_F2I);
+ } else if (DestRC == &NVPTX::Int128RegsRegClass) {
+ Op = NVPTX::IMOV128rr;
} else if (DestRC == &NVPTX::Float32RegsRegClass) {
Op = (SrcRC == &NVPTX::Float32RegsRegClass ? NVPTX::FMOV32rr
: NVPTX::BITCONVERT_32_I2F);
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index c4c35a1f74ba9..827febe845a4c 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -2097,6 +2097,8 @@ let IsSimpleMove=1, hasSideEffects=0 in {
"mov.u32 \t$dst, $sss;", []>;
def IMOV64rr : NVPTXInst<(outs Int64Regs:$dst), (ins Int64Regs:$sss),
"mov.u64 \t$dst, $sss;", []>;
+ def IMOV128rr : NVPTXInst<(outs Int128Regs:$dst), (ins Int128Regs:$sss),
+ "mov.b128 \t$dst, $sss;", []>;
def IMOVB16rr : NVPTXInst<(outs Int16Regs:$dst), (ins Int16Regs:$sss),
"mov.b16 \t$dst, $sss;", []>;
@@ -3545,6 +3547,9 @@ let hasSideEffects = false in {
def V2I32toI64 : NVPTXInst<(outs Int64Regs:$d),
(ins Int32Regs:$s1, Int32Regs:$s2),
"mov.b64 \t$d, {{$s1, $s2}};", []>;
+ def V2I64toI128 : NVPTXInst<(outs Int128Regs:$d),
+ (ins Int64Regs:$s1, Int64Regs:$s2),
+ "mov.b128 \t$d, {{$s1, $s2}};", []>;
def V2F32toF64 : NVPTXInst<(outs Float64Regs:$d),
(ins Float32Regs:$s1, Float32Regs:$s2),
"mov.b64 \t$d, {{$s1, $s2}};", []>;
@@ -3560,6 +3565,9 @@ let hasSideEffects = false in {
def I64toV2I32 : NVPTXInst<(outs Int32Regs:$d1, Int32Regs:$d2),
(ins Int64Regs:$s),
"mov.b64 \t{{$d1, $d2}}, $s;", []>;
+ def I128toV2I64: NVPTXInst<(outs Int64Regs:$d1, Int64Regs:$d2),
+ (ins Int128Regs:$s),
+ "mov.b128 \t{{$d1, $d2}}, $s;", []>;
def F64toV2F32 : NVPTXInst<(outs Float32Regs:$d1, Float32Regs:$d2),
(ins Float64Regs:$s),
"mov.b64 \t{{$d1, $d2}}, $s;", []>;
@@ -3629,7 +3637,7 @@ def : Pat<(i32 (ctlz (i32 Int32Regs:$a))), (CLZr32 Int32Regs:$a)>;
// ptx value to 64 bits to match the ISD node's semantics, unless we know we're
// truncating back down to 32 bits.
def : Pat<(i64 (ctlz Int64Regs:$a)), (CVT_u64_u32 (CLZr64 Int64Regs:$a), CvtNONE)>;
-def : Pat<(i32 (trunc (ctlz Int64Regs:$a))), (CLZr64 Int64Regs:$a)>;
+def : Pat<(i32 (trunc (i64 (ctlz Int64Regs:$a)))), (CLZr64 Int64Regs:$a)>;
// For 16-bit ctlz, we zero-extend to 32-bit, perform the count, then trunc the
// result back to 16-bits if necessary. We also need to subtract 16 because
@@ -3667,7 +3675,7 @@ def : Pat<(i32 (ctpop (i32 Int32Regs:$a))), (POPCr32 Int32Regs:$a)>;
// pattern that avoids the type conversion if we're truncating the result to
// i32 anyway.
def : Pat<(ctpop Int64Regs:$a), (CVT_u64_u32 (POPCr64 Int64Regs:$a), CvtNONE)>;
-def : Pat<(i32 (trunc (ctpop Int64Regs:$a))), (POPCr64 Int64Regs:$a)>;
+def : Pat<(i32 (trunc (i64 (ctpop Int64Regs:$a)))), (POPCr64 Int64Regs:$a)>;
// For 16-bit, we zero-extend to 32-bit, then trunc the result back to 16-bits.
// If we know that we're storing into an i32, we can avoid the final trunc.
diff --git a/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp b/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
index f1213f030bba7..a8a23f04c1249 100644
--- a/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
@@ -31,6 +31,8 @@ std::string getNVPTXRegClassName(TargetRegisterClass const *RC) {
return ".f32";
if (RC == &NVPTX::Float64RegsRegClass)
return ".f64";
+ if (RC == &NVPTX::Int128RegsRegClass)
+ return ".b128";
if (RC == &NVPTX::Int64RegsRegClass)
// We use untyped (.b) integer registers here as NVCC does.
// Correctness of generated code does not depend on register type,
@@ -67,6 +69,8 @@ std::string getNVPTXRegClassStr(TargetRegisterClass const *RC) {
return "%f";
if (RC == &NVPTX::Float64RegsRegClass)
return "%fd";
+ if (RC == &NVPTX::Int128RegsRegClass)
+ return "%rq";
if (RC == &NVPTX::Int64RegsRegClass)
return "%rd";
if (RC == &NVPTX::Int32RegsRegClass)
diff --git a/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.td b/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.td
index b5231a9cf67f9..2011f0f7e328f 100644
--- a/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.td
@@ -37,6 +37,7 @@ foreach i = 0...4 in {
def RS#i : NVPTXReg<"%rs"#i>; // 16-bit
def R#i : NVPTXReg<"%r"#i>; // 32-bit
def RL#i : NVPTXReg<"%rd"#i>; // 64-bit
+ def RQ#i : NVPTXReg<"%rq"#i>; // 128-bit
def H#i : NVPTXReg<"%h"#i>; // 16-bit float
def HH#i : NVPTXReg<"%hh"#i>; // 2x16-bit float
def F#i : NVPTXReg<"%f"#i>; // 32-bit float
@@ -62,6 +63,8 @@ def Int32Regs : NVPTXRegClass<[i32, v2f16, v2bf16, v2i16, v4i8], 32,
(add (sequence "R%u", 0, 4),
VRFrame32, VRFrameLocal32)>;
def Int64Regs : NVPTXRegClass<[i64], 64, (add (sequence "RL%u", 0, 4), VRFrame64, VRFrameLocal64)>;
+// 128-bit regs are not defined as general regs in NVPTX. They are used for inlineASM only.
+def Int128Regs : NVPTXRegClass<[i128], 128, (add (sequence "RQ%u", 0, 4))>;
def Float32Regs : NVPTXRegClass<[f32], 32, (add (sequence "F%u", 0, 4))>;
def Float64Regs : NVPTXRegClass<[f64], 64, (add (sequence "FL%u", 0, 4))>;
def Int32ArgRegs : NVPTXRegClass<[i32], 32, (add (sequence "ia%u", 0, 4))>;
diff --git a/llvm/test/CodeGen/NVPTX/inline-asm-b128-test1.ll b/llvm/test/CodeGen/NVPTX/inline-asm-b128-test1.ll
new file mode 100644
index 0000000000000..dec0451c34ccc
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/inline-asm-b128-test1.ll
@@ -0,0 +1,92 @@
+; RUN: llc < %s -march=nvptx -mcpu=sm_70 -o - 2>&1 | FileCheck %s
+
+target triple = "nvptx64-nvidia-cuda"
+
+ at value = internal addrspace(1) global i128 0, align 16
+ at llvm.used = appending global [6 x ptr] [ptr @_Z7kernel1v, ptr @_Z7kernel2Pn, ptr @_Z7kernel3Pb, ptr @_Z7kernel4v, ptr @_Z7kernel5Pn, ptr addrspacecast (ptr addrspace(1) @value to ptr)], section "llvm.metadata"
+
+; Function Attrs: alwaysinline convergent mustprogress willreturn
+define void @_Z7kernel1v() #0 {
+ ; CHECK-LABEL: _Z7kernel1v
+ ; CHECK: mov.u64 [[REG_HI:%rd[0-9]+]], 0;
+ ; CHECK: mov.u64 [[REG_LO:%rd[0-9]+]], 42;
+ ; CHECK: mov.b128 [[REG_128:%rq[0-9]+]], {[[REG_LO]], [[REG_HI]]};
+ ; CHECK: { st.b128 [{{%rd[0-9]+}}], [[REG_128]]; }
+
+ tail call void asm sideeffect "{ st.b128 [$0], $1; }", "l,q"(ptr nonnull addrspacecast (ptr addrspace(1) @value to ptr), i128 42) #3
+ ret void
+}
+
+; Function Attrs: alwaysinline convergent mustprogress willreturn
+define void @_Z7kernel2Pn(ptr nocapture readonly %data) #0 {
+ ; CHECK-LABEL: _Z7kernel2Pn
+ ; CHECK: ld.global.u64 [[REG_HI:%rd[0-9]+]], [[[REG_Addr:%r[0-9]+]]+8];
+ ; CHECK: ld.global.u64 [[REG_LO:%rd[0-9]+]], [[[REG_Addr]]];
+ ; CHECK: mov.b128 [[REG_128:%rq[0-9]+]], {[[REG_LO]], [[REG_HI]]};
+ ; CHECK: { st.b128 [{{%rd[0-9]+}}], [[REG_128]]; }
+
+ %1 = addrspacecast ptr %data to ptr addrspace(1)
+ %2 = load <2 x i64>, ptr addrspace(1) %1, align 16
+ %3 = bitcast <2 x i64> %2 to i128
+ tail call void asm sideeffect "{ st.b128 [$0], $1; }", "l,q"(ptr nonnull addrspacecast (ptr addrspace(1) @value to ptr), i128 %3) #3
+ ret void
+}
+
+; Function Attrs: alwaysinline convergent mustprogress willreturn
+define void @_Z7kernel3Pb(ptr nocapture readonly %flag) #0 {
+ ; CHECK-LABEL: _Z7kernel3Pb
+ ; CHECK: selp.b64 [[REG_LO:%rd[0-9]+]], 24, 42, {{%p[0-9]+}};
+ ; CHECK: mov.u64 [[REG_HI:%rd[0-9]+]], 0;
+ ; CHECK: mov.b128 [[REG_128:%rq[0-9]+]], {[[REG_LO]], [[REG_HI]]};
+ ; CHECK: { st.b128 [{{%rd[0-9]+}}], [[REG_128]]; }
+
+ %1 = addrspacecast ptr %flag to ptr addrspace(1)
+ %tmp1 = load i8, ptr addrspace(1) %1, align 1
+ %tobool.not = icmp eq i8 %tmp1, 0
+ %. = select i1 %tobool.not, i128 24, i128 42
+ tail call void asm sideeffect "{ st.b128 [$0], $1; }", "l,q"(ptr nonnull addrspacecast (ptr addrspace(1) @value to ptr), i128 %.) #3
+ ret void
+}
+
+; Function Attrs: alwaysinline mustprogress willreturn memory(write, argmem: none, inaccessiblemem: none)
+define void @_Z7kernel4v() #1 {
+ ; CHECK-LABEL: _Z7kernel4v
+ ; CHECK-O3: { mov.b128 [[REG_128:%rq[0-9]+]], 41; }
+ ; CHECK-O3: mov.b128 {%rd{{[0-9]+}}, %rd{{[0-9]+}}}, [[REG_128]];
+
+ %1 = tail call i128 asm "{ mov.b128 $0, 41; }", "=q"() #4
+ %add = add nsw i128 %1, 1
+ %2 = bitcast i128 %add to <2 x i64>
+ store <2 x i64> %2, ptr addrspace(1) @value, align 16
+ ret void
+}
+
+; Function Attrs: alwaysinline mustprogress willreturn memory(write, argmem: read, inaccessiblemem: none)
+define void @_Z7kernel5Pn(ptr nocapture readonly %data) #2 {
+ ; CHECK-LABEL: _Z7kernel5Pn
+ ; CHECK-O3: ld.global.v2.u64 {[[REG_LO_IN:%rd[0-9]+]], [[REG_HI_IN:%rd[0-9]+]]}, [{{%rd[0-9]+}}];
+ ; CHECK-O3: mov.b128 [[REG_128_IN:%rq[0-9]+]], {[[REG_LO_IN]], [[REG_HI_IN]]};
+ ; CHECK-O3: { mov.b128 [[REG_128_OUT:%rq[0-9]+]], [[REG_128_IN]]; }
+ ; CHECK-O3: mov.b128 {%rd{{[0-9]+}}, %rd{{[0-9]+}}}, [[REG_128_OUT]];
+
+ %1 = addrspacecast ptr %data to ptr addrspace(1)
+ %2 = load <2 x i64>, ptr addrspace(1) %1, align 16
+ %3 = bitcast <2 x i64> %2 to i128
+ %4 = tail call i128 asm "{ mov.b128 $0, $1; }", "=q,q"(i128 %3) #4
+ %add = add nsw i128 %4, 1
+ %5 = bitcast i128 %add to <2 x i64>
+ store <2 x i64> %5, ptr addrspace(1) @value, align 16
+ ret void
+}
+
+attributes #0 = { alwaysinline convergent mustprogress willreturn "nvvm.annotations_transplanted" "nvvm.kernel" "nvvm.restrict_processed" "target-cpu"="sm_89" }
+attributes #1 = { alwaysinline mustprogress willreturn memory(write, argmem: none, inaccessiblemem: none) "nvvm.annotations_transplanted" "nvvm.kernel" "nvvm.restrict_processed" "target-cpu"="sm_89" }
+attributes #2 = { alwaysinline mustprogress willreturn memory(write, argmem: read, inaccessiblemem: none) "nvvm.annotations_transplanted" "nvvm.kernel" "nvvm.restrict_processed" "target-cpu"="sm_89" }
+attributes #3 = { convergent nounwind }
+attributes #4 = { nounwind }
+
+
+!nvvmir.version = !{!0, !1, !0, !1, !1, !0, !0, !0, !1}
+
+!0 = !{i32 2, i32 0, i32 3, i32 1}
+!1 = !{i32 2, i32 0}
diff --git a/llvm/test/CodeGen/NVPTX/inline-asm-b128-test2.ll b/llvm/test/CodeGen/NVPTX/inline-asm-b128-test2.ll
new file mode 100644
index 0000000000000..337479a06c3f0
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/inline-asm-b128-test2.ll
@@ -0,0 +1,57 @@
+; RUN: llc < %s -march=nvptx -mcpu=sm_70 -o - 2>&1 | FileCheck %s
+
+target triple = "nvptx64-nvidia-cuda"
+
+ at u128_max = internal addrspace(1) global i128 0, align 16
+ at u128_zero = internal addrspace(1) global i128 0, align 16
+ at i128_max = internal addrspace(1) global i128 0, align 16
+ at i128_min = internal addrspace(1) global i128 0, align 16
+ at v_u128_max = internal addrspace(1) global i128 0, align 16
+ at v_u128_zero = internal addrspace(1) global i128 0, align 16
+ at v_i128_max = internal addrspace(1) global i128 0, align 16
+ at v_i128_min = internal addrspace(1) global i128 0, align 16
+ at v64 = internal addrspace(1) global i64* null, align 8
+ at llvm.used = appending global [10 x i8*] [i8* bitcast (i128* addrspacecast (i128 addrspace(1)* @u128_max to i128*) to i8*), i8* bitcast (i128* addrspacecast (i128 addrspace(1)* @u128_zero to i128*) to i8*), i8* bitcast (i128* addrspacecast (i128 addrspace(1)* @i128_max to i128*) to i8*), i8* bitcast (i128* addrspacecast (i128 addrspace(1)* @i128_min to i128*) to i8*), i8* bitcast (i128* addrspacecast (i128 addrspace(1)* @v_u128_max to i128*) to i8*), i8* bitcast (i128* addrspacecast (i128 addrspace(1)* @v_u128_zero to i128*) to i8*), i8* bitcast (i128* addrspacecast (i128 addrspace(1)* @v_i128_max to i128*) to i8*), i8* bitcast (i128* addrspacecast (i128 addrspace(1)* @v_i128_min to i128*) to i8*), i8* bitcast (i64** addrspacecast (i64* addrspace(1)* @v64 to i64**) to i8*), i8* bitcast (void ()* @_Z6kernelv to i8*)], section "llvm.metadata"
+
+; Function Attrs: alwaysinline
+define void @_Z6kernelv() #0 {
+ ; CHECK-LABLE: _Z6kernelv
+ ; CHECK: mov.u64 [[U64_MAX:%rd[0-9]+]], -1;
+ ; CHECK: mov.b128 [[U128_MAX:%rq[0-9]+]], {[[U64_MAX]], [[U64_MAX]]};
+ ; CHECK: mov.u64 [[I128_MAX_HI:%rd[0-9]+]], 9223372036854775807;
+ ; CHECK: mov.b128 [[I128_MAX:%rq[0-9]+]], {[[U64_MAX]], [[I128_MAX_HI]]};
+ ; CHECK: mov.u64 [[I128_MIN_HI:%rd[0-9]+]], -9223372036854775808;
+ ; CHECK: mov.u64 [[ZERO:%rd[0-9]+]], 0;
+ ; CHECK: mov.b128 [[I128_MIN:%rq[0-9]+]], {[[ZERO]], [[I128_MIN_HI]]};
+ ; CHECK: mov.b128 [[U128_ZERO:%rq[0-9]+]], {[[ZERO]], [[ZERO]]};
+
+ %tmp = load i64*, i64** addrspacecast (i64* addrspace(1)* @v64 to i64**), align 8
+ %add.ptr = getelementptr inbounds i64, i64* %tmp, i32 0
+ %tmp1 = load i64*, i64** addrspacecast (i64* addrspace(1)* @v64 to i64**), align 8
+ %add.ptr2 = getelementptr inbounds i64, i64* %tmp1, i32 1
+ call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}\0A\09", "q,l,l,l"(i128 -1, i64* %add.ptr, i64* %add.ptr2, i128* addrspacecast (i128 addrspace(1)* @v_u128_max to i128*)) #1
+ %tmp3 = load i64*, i64** addrspacecast (i64* addrspace(1)* @v64 to i64**), align 8
+ %add.ptr4 = getelementptr inbounds i64, i64* %tmp3, i32 2
+ %tmp5 = load i64*, i64** addrspacecast (i64* addrspace(1)* @v64 to i64**), align 8
+ %add.ptr6 = getelementptr inbounds i64, i64* %tmp5, i32 3
+ call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}\0A\09", "q,l,l,l"(i128 170141183460469231731687303715884105727, i64* %add.ptr4, i64* %add.ptr6, i128* addrspacecast (i128 addrspace(1)* @v_i128_max to i128*)) #1
+ %tmp7 = load i64*, i64** addrspacecast (i64* addrspace(1)* @v64 to i64**), align 8
+ %add.ptr8 = getelementptr inbounds i64, i64* %tmp7, i32 4
+ %tmp9 = load i64*, i64** addrspacecast (i64* addrspace(1)* @v64 to i64**), align 8
+ %add.ptr10 = getelementptr inbounds i64, i64* %tmp9, i32 5
+ call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}\0A\09", "q,l,l,l"(i128 -170141183460469231731687303715884105728, i64* %add.ptr8, i64* %add.ptr10, i128* addrspacecast (i128 addrspace(1)* @v_i128_min to i128*)) #1
+ %tmp11 = load i64*, i64** addrspacecast (i64* addrspace(1)* @v64 to i64**), align 8
+ %add.ptr12 = getelementptr inbounds i64, i64* %tmp11, i32 6
+ %tmp13 = load i64*, i64** addrspacecast (i64* addrspace(1)* @v64 to i64**), align 8
+ %add.ptr14 = getelementptr inbounds i64, i64* %tmp13, i32 7
+ call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}\0A\09", "q,l,l,l"(i128 0, i64* %add.ptr12, i64* %add.ptr14, i128* addrspacecast (i128 addrspace(1)* @v_u128_zero to i128*)) #1
+ ret void
+}
+
+attributes #0 = { alwaysinline "nvvm.annotations_transplanted" "nvvm.kernel" }
+attributes #1 = { nounwind }
+
+!nvvmir.version = !{!0, !1, !0, !1, !1, !0, !0, !0, !1}
+
+!0 = !{i32 2, i32 0, i32 3, i32 1}
+!1 = !{i32 2, i32 0}
diff --git a/llvm/test/CodeGen/NVPTX/inline-asm-b128-test3.ll b/llvm/test/CodeGen/NVPTX/inline-asm-b128-test3.ll
new file mode 100644
index 0000000000000..4f077ec5383c9
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/inline-asm-b128-test3.ll
@@ -0,0 +1,179 @@
+; RUN: llc < %s -march=nvptx -mcpu=sm_70 -o - 2>&1 | FileCheck %s
+
+target triple = "nvptx64-nvidia-cuda"
+
+ at size = internal addrspace(1) global i32 0, align 4
+ at value = internal addrspace(1) global i128 0, align 16
+ at x = internal addrspace(1) global i128 0, align 16
+ at y = internal addrspace(1) global i128 0, align 16
+ at z = internal addrspace(1) global i128 0, align 16
+ at llvm.used = appending global [6 x ptr] [ptr @_Z6kernelv, ptr addrspacecast (ptr addrspace(1) @size to ptr), ptr addrspacecast (ptr addrspace(1) @value to ptr), ptr addrspacecast (ptr addrspace(1) @x to ptr), ptr addrspacecast (ptr addrspace(1) @y to ptr), ptr addrspacecast (ptr addrspace(1) @z to ptr)], section "llvm.metadata"
+
+; Function Attrs: alwaysinline mustprogress willreturn memory(readwrite, argmem: none, inaccessiblemem: none)
+define void @_Z6kernelv() #0 {
+ ; CHECK-LABEL: _Z6kernelv
+ ; CHECK: mov.b128 [[X:%rq[0-9]+]], {{{%rd[0-9]+}}, {{%rd[0-9]+}}};
+ ; CHECK: mov.b128 [[Y:%rq[0-9]+]], {{{%rd[0-9]+}}, {{%rd[0-9]+}}};
+ ; CHECK: mov.b128 [[Z:%rq[0-9]+]], {{{%rd[0-9]+}}, {{%rd[0-9]+}}};
+ ; CHECK: mov.b128 {lo, hi}, [[X]];
+ ; CHECK: mov.b128 [[X]], {lo, hi};
+ ; CHECK: mov.b128 {lo, hi}, [[Y]];
+ ; CHECK: mov.b128 [[Y]], {lo, hi};
+ ; CHECK: mov.b128 {lo, hi}, [[Z]];
+ ; CHECK: mov.b128 [[Z]], {lo, hi};
+ ; CHECK: mov.b128 {[[X_LO:%rd[0-9]+]], [[X_HI:%rd[0-9]+]]}, [[X]];
+ ; CHECK: mov.b128 {[[Y_LO:%rd[0-9]+]], [[Y_HI:%rd[0-9]+]]}, [[Y]];
+ ; CHECK: mov.b128 {[[Z_LO:%rd[0-9]+]], [[Z_HI:%rd[0-9]+]]}, [[Z]];
+ ; CHECK: mov.b128 [[X_NEW:%rq[0-9]+]], {[[X_LO]], [[X_HI]]};
+ ; CHECK: mov.b128 [[Y_NEW:%rq[0-9]+]], {[[Y_LO]], [[Y_HI]]};
+ ; CHECK: mov.b128 [[Z_NEW:%rq[0-9]+]], {[[Z_LO]], [[Z_HI]]};
+ ; CHECK: mov.b128 {lo, hi}, [[X_NEW]];
+ ; CHECK: mov.b128 [[X_NEW]], {lo, hi};
+ ; CHECK: mov.b128 {lo, hi}, [[Y_NEW]];
+ ; CHECK: mov.b128 [[Y_NEW]], {lo, hi};
+ ; CHECK: mov.b128 {lo, hi}, [[Z_NEW]];
+ ; CHECK: mov.b128 [[Z_NEW]], {lo, hi};
+ ; CHECK: mov.b128 {{{%rd[0-9]+}}, {{%rd[0-9]+}}}, [[X_NEW]];
+ ; CHECK: mov.b128 {{{%rd[0-9]+}}, {{%rd[0-9]+}}}, [[Y_NEW]];
+ ; CHECK: mov.b128 {{{%rd[0-9]+}}, {{%rd[0-9]+}}}, [[Z_NEW]];
+
+ %tmp11 = load i32, ptr addrspace(1) @size, align 4
+ %cmp3.not = icmp eq i32 %tmp11, 0
+ br i1 %cmp3.not, label %._crit_edge, label %.lr.ph.preheader
+
+.lr.ph.preheader: ; preds = %0
+ %x.promoted5 = load i128, ptr addrspace(1) @x, align 16
+ %y.promoted6 = load i128, ptr addrspace(1) @y, align 16
+ %z.promoted7 = load i128, ptr addrspace(1) @z, align 16
+ %value.promoted8 = load i128, ptr addrspace(1) @value, align 16
+ %umax = sext i32 %tmp11 to i64
+ %xtraiter = and i64 %umax, 3
+ %1 = icmp ult i32 %tmp11, 4
+ br i1 %1, label %._crit_edge.loopexit.unr-lcssa, label %.lr.ph.preheader.new
+
+.lr.ph.preheader.new: ; preds = %.lr.ph.preheader
+ %unroll_iter = and i64 %umax, -4
+ br label %.lr.ph
+
+.lr.ph: ; preds = %.lr.ph, %.lr.ph.preheader.new
+ %2 = phi i128 [ %value.promoted8, %.lr.ph.preheader.new ], [ %add14.3, %.lr.ph ]
+ %3 = phi i128 [ %z.promoted7, %.lr.ph.preheader.new ], [ %asmresult21.3, %.lr.ph ]
+ %4 = phi i128 [ %y.promoted6, %.lr.ph.preheader.new ], [ %asmresult20.3, %.lr.ph ]
+ %5 = phi i128 [ %x.promoted5, %.lr.ph.preheader.new ], [ %asmresult19.3, %.lr.ph ]
+ %i.04 = phi i64 [ 0, %.lr.ph.preheader.new ], [ %inc.3, %.lr.ph ]
+ %niter = phi i64 [ 0, %.lr.ph.preheader.new ], [ %niter.next.3, %.lr.ph ]
+ %6 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09add.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09add.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09add.cc.u64 lo, lo, 3;\0A\09add.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %i.04, i128 %5, i128 %4, i128 %3) #1
+ %asmresult = extractvalue { i128, i128, i128 } %6, 0
+ %asmresult7 = extractvalue { i128, i128, i128 } %6, 1
+ %asmresult8 = extractvalue { i128, i128, i128 } %6, 2
+ %add = add nsw i128 %asmresult, %asmresult7
+ %add12 = add nsw i128 %add, %asmresult8
+ %add14 = add nsw i128 %add12, %2
+ %7 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09sub.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09sub.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09sub.cc.u64 lo, lo, 3;\0A\09sub.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %i.04, i128 %asmresult, i128 %asmresult7, i128 %asmresult8) #1
+ %asmresult19 = extractvalue { i128, i128, i128 } %7, 0
+ %asmresult20 = extractvalue { i128, i128, i128 } %7, 1
+ %asmresult21 = extractvalue { i128, i128, i128 } %7, 2
+ %inc = add nuw nsw i64 %i.04, 1
+ %8 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09add.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09add.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09add.cc.u64 lo, lo, 3;\0A\09add.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %inc, i128 %asmresult19, i128 %asmresult20, i128 %asmresult21) #1
+ %asmresult.1 = extractvalue { i128, i128, i128 } %8, 0
+ %asmresult7.1 = extractvalue { i128, i128, i128 } %8, 1
+ %asmresult8.1 = extractvalue { i128, i128, i128 } %8, 2
+ %add.1 = add nsw i128 %asmresult.1, %asmresult7.1
+ %add12.1 = add nsw i128 %add.1, %asmresult8.1
+ %add14.1 = add nsw i128 %add12.1, %add14
+ %9 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09sub.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09sub.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09sub.cc.u64 lo, lo, 3;\0A\09sub.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %inc, i128 %asmresult.1, i128 %asmresult7.1, i128 %asmresult8.1) #1
+ %asmresult19.1 = extractvalue { i128, i128, i128 } %9, 0
+ %asmresult20.1 = extractvalue { i128, i128, i128 } %9, 1
+ %asmresult21.1 = extractvalue { i128, i128, i128 } %9, 2
+ %inc.1 = add nuw nsw i64 %i.04, 2
+ %10 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09add.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09add.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09add.cc.u64 lo, lo, 3;\0A\09add.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %inc.1, i128 %asmresult19.1, i128 %asmresult20.1, i128 %asmresult21.1) #1
+ %asmresult.2 = extractvalue { i128, i128, i128 } %10, 0
+ %asmresult7.2 = extractvalue { i128, i128, i128 } %10, 1
+ %asmresult8.2 = extractvalue { i128, i128, i128 } %10, 2
+ %add.2 = add nsw i128 %asmresult.2, %asmresult7.2
+ %add12.2 = add nsw i128 %add.2, %asmresult8.2
+ %add14.2 = add nsw i128 %add12.2, %add14.1
+ %11 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09sub.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09sub.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09sub.cc.u64 lo, lo, 3;\0A\09sub.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %inc.1, i128 %asmresult.2, i128 %asmresult7.2, i128 %asmresult8.2) #1
+ %asmresult19.2 = extractvalue { i128, i128, i128 } %11, 0
+ %asmresult20.2 = extractvalue { i128, i128, i128 } %11, 1
+ %asmresult21.2 = extractvalue { i128, i128, i128 } %11, 2
+ %inc.2 = add nuw nsw i64 %i.04, 3
+ %12 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09add.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09add.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09add.cc.u64 lo, lo, 3;\0A\09add.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %inc.2, i128 %asmresult19.2, i128 %asmresult20.2, i128 %asmresult21.2) #1
+ %asmresult.3 = extractvalue { i128, i128, i128 } %12, 0
+ %asmresult7.3 = extractvalue { i128, i128, i128 } %12, 1
+ %asmresult8.3 = extractvalue { i128, i128, i128 } %12, 2
+ %add.3 = add nsw i128 %asmresult.3, %asmresult7.3
+ %add12.3 = add nsw i128 %add.3, %asmresult8.3
+ %add14.3 = add nsw i128 %add12.3, %add14.2
+ %13 = bitcast i128 %add14.3 to <2 x i64>
+ store <2 x i64> %13, ptr addrspace(1) @value, align 16
+ %14 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09sub.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09sub.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09sub.cc.u64 lo, lo, 3;\0A\09sub.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %inc.2, i128 %asmresult.3, i128 %asmresult7.3, i128 %asmresult8.3) #1
+ %asmresult19.3 = extractvalue { i128, i128, i128 } %14, 0
+ %asmresult20.3 = extractvalue { i128, i128, i128 } %14, 1
+ %asmresult21.3 = extractvalue { i128, i128, i128 } %14, 2
+ %15 = bitcast i128 %asmresult19.3 to <2 x i64>
+ store <2 x i64> %15, ptr addrspace(1) @x, align 16
+ %16 = bitcast i128 %asmresult20.3 to <2 x i64>
+ store <2 x i64> %16, ptr addrspace(1) @y, align 16
+ %17 = bitcast i128 %asmresult21.3 to <2 x i64>
+ store <2 x i64> %17, ptr addrspace(1) @z, align 16
+ %inc.3 = add nuw i64 %i.04, 4
+ %niter.next.3 = add i64 %niter, 4
+ %niter.ncmp.3.not = icmp eq i64 %niter.next.3, %unroll_iter
+ br i1 %niter.ncmp.3.not, label %._crit_edge.loopexit.unr-lcssa, label %.lr.ph, !llvm.loop !2
+
+._crit_edge.loopexit.unr-lcssa: ; preds = %.lr.ph, %.lr.ph.preheader
+ %.unr = phi i128 [ %value.promoted8, %.lr.ph.preheader ], [ %add14.3, %.lr.ph ]
+ %.unr9 = phi i128 [ %z.promoted7, %.lr.ph.preheader ], [ %asmresult21.3, %.lr.ph ]
+ %.unr10 = phi i128 [ %y.promoted6, %.lr.ph.preheader ], [ %asmresult20.3, %.lr.ph ]
+ %.unr11 = phi i128 [ %x.promoted5, %.lr.ph.preheader ], [ %asmresult19.3, %.lr.ph ]
+ %i.04.unr = phi i64 [ 0, %.lr.ph.preheader ], [ %inc.3, %.lr.ph ]
+ %lcmp.mod.not = icmp eq i64 %xtraiter, 0
+ br i1 %lcmp.mod.not, label %._crit_edge, label %.lr.ph.epil
+
+.lr.ph.epil: ; preds = %.lr.ph.epil, %._crit_edge.loopexit.unr-lcssa
+ %18 = phi i128 [ %add14.epil, %.lr.ph.epil ], [ %.unr, %._crit_edge.loopexit.unr-lcssa ]
+ %19 = phi i128 [ %asmresult21.epil, %.lr.ph.epil ], [ %.unr9, %._crit_edge.loopexit.unr-lcssa ]
+ %20 = phi i128 [ %asmresult20.epil, %.lr.ph.epil ], [ %.unr10, %._crit_edge.loopexit.unr-lcssa ]
+ %21 = phi i128 [ %asmresult19.epil, %.lr.ph.epil ], [ %.unr11, %._crit_edge.loopexit.unr-lcssa ]
+ %i.04.epil = phi i64 [ %inc.epil, %.lr.ph.epil ], [ %i.04.unr, %._crit_edge.loopexit.unr-lcssa ]
+ %epil.iter = phi i64 [ %epil.iter.next, %.lr.ph.epil ], [ 0, %._crit_edge.loopexit.unr-lcssa ]
+ %22 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09add.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09add.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09add.cc.u64 lo, lo, 3;\0A\09add.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %i.04.epil, i128 %21, i128 %20, i128 %19) #1
+ %asmresult.epil = extractvalue { i128, i128, i128 } %22, 0
+ %asmresult7.epil = extractvalue { i128, i128, i128 } %22, 1
+ %asmresult8.epil = extractvalue { i128, i128, i128 } %22, 2
+ %add.epil = add nsw i128 %asmresult.epil, %asmresult7.epil
+ %add12.epil = add nsw i128 %add.epil, %asmresult8.epil
+ %add14.epil = add nsw i128 %add12.epil, %18
+ %23 = bitcast i128 %add14.epil to <2 x i64>
+ store <2 x i64> %23, ptr addrspace(1) @value, align 16
+ %24 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09sub.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09sub.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09sub.cc.u64 lo, lo, 3;\0A\09sub.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %i.04.epil, i128 %asmresult.epil, i128 %asmresult7.epil, i128 %asmresult8.epil) #1
+ %asmresult19.epil = extractvalue { i128, i128, i128 } %24, 0
+ %asmresult20.epil = extractvalue { i128, i128, i128 } %24, 1
+ %asmresult21.epil = extractvalue { i128, i128, i128 } %24, 2
+ %25 = bitcast i128 %asmresult19.epil to <2 x i64>
+ store <2 x i64> %25, ptr addrspace(1) @x, align 16
+ %26 = bitcast i128 %asmresult20.epil to <2 x i64>
+ store <2 x i64> %26, ptr addrspace(1) @y, align 16
+ %27 = bitcast i128 %asmresult21.epil to <2 x i64>
+ store <2 x i64> %27, ptr addrspace(1) @z, align 16
+ %inc.epil = add nuw i64 %i.04.epil, 1
+ %epil.iter.next = add i64 %epil.iter, 1
+ %epil.iter.cmp.not = icmp eq i64 %epil.iter.next, %xtraiter
+ br i1 %epil.iter.cmp.not, label %._crit_edge, label %.lr.ph.epil, !llvm.loop !4
+
+._crit_edge: ; preds = %.lr.ph.epil, %._crit_edge.loopexit.unr-lcssa, %0
+ ret void
+}
+
+attributes #0 = { alwaysinline mustprogress willreturn memory(readwrite, argmem: none, inaccessiblemem: none) "nvvm.annotations_transplanted" "nvvm.kernel" "nvvm.restrict_processed" "target-cpu"="sm_89" }
+attributes #1 = { nounwind }
+
+!nvvmir.version = !{!0, !1, !0, !1, !1, !0, !0, !0, !1}
+
+!0 = !{i32 2, i32 0, i32 3, i32 1}
+!1 = !{i32 2, i32 0}
+!2 = distinct !{!2, !3}
+!3 = !{!"llvm.loop.mustprogress"}
+!4 = distinct !{!4, !5}
+!5 = !{!"llvm.loop.unroll.disable"}
>From d3d53434eda400a44121c4a2e52dcaa8b9077227 Mon Sep 17 00:00:00 2001
From: chengjunp <chengjunp at nvidia.com>
Date: Tue, 25 Jun 2024 21:54:48 +0000
Subject: [PATCH 2/6] Format Code & Update tests
---
llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp | 17 ++---
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 23 +++---
.../CodeGen/NVPTX/inline-asm-b128-test1.ll | 28 +++-----
.../CodeGen/NVPTX/inline-asm-b128-test2.ll | 70 ++++++++-----------
.../CodeGen/NVPTX/inline-asm-b128-test3.ll | 24 +++----
5 files changed, 73 insertions(+), 89 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
index 05706e200bda6..9c0498560db21 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
@@ -527,7 +527,7 @@ void NVPTXDAGToDAGISel::Select(SDNode *N) {
break;
}
case ISD::CopyFromReg: {
- if(N->getOperand(1).getValueType() == MVT::i128){
+ if (N->getOperand(1).getValueType() == MVT::i128) {
SelectI128toV2I64(N);
return;
}
@@ -3825,17 +3825,17 @@ void NVPTXDAGToDAGISel::SelectV2I64toI128(SDNode *N) {
SDValue Dst = N->getOperand(1);
SDValue Lo = N->getOperand(2);
SDValue Hi = N->getOperand(3);
-
+
SDLoc DL(N);
SDNode *Mov =
CurDAG->getMachineNode(NVPTX::V2I64toI128, DL, MVT::i128, {Lo, Hi});
-
+
SmallVector<EVT, 8> ResultsType(N->value_begin(), N->value_end());
SmallVector<SDValue, 8> NewOps(N->getNumOperands() - 1);
NewOps[0] = N->getOperand(0);
NewOps[1] = Dst;
NewOps[2] = SDValue(Mov, 0);
- if (N->getNumOperands() == 5)
+ if (N->getNumOperands() == 5)
NewOps[3] = N->getOperand(4);
SDValue NewValue = CurDAG->getNode(ISD::CopyToReg, DL, ResultsType, NewOps);
@@ -3847,17 +3847,18 @@ void NVPTXDAGToDAGISel::SelectI128toV2I64(SDNode *N) {
// Dst:i128, Src:i128
//
// {lo, hi} = CopyFromReg Src
- //
+ //
// ==>
- //
+ //
// {lo, hi} = I128toV2I64 Src
- //
+ //
SDValue Ch = N->getOperand(0);
SDValue Src = N->getOperand(1);
SDValue Glue = N->getOperand(2);
SDLoc DL(N);
- // Add Glue and Ch to the operands and results to avoid break the execution order
+ // Add Glue and Ch to the operands and results to avoid break the execution
+ // order
SDNode *Mov = CurDAG->getMachineNode(
NVPTX::I128toV2I64, DL,
{MVT::i64, MVT::i64, Ch.getValueType(), Glue.getValueType()},
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 927c3ef55dcb1..3cfac3baeb5c9 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -3107,15 +3107,15 @@ SDValue NVPTXTargetLowering::LowerCopyToReg_128(SDValue Op,
assert(Op.getOperand(1).getValueType() == MVT::i128 &&
"Custom lowering for 128-bit CopyToReg only");
-
+
SDNode *Node = Op.getNode();
SDLoc DL(Node);
SDValue Cast = DAG.getNode(ISD::BITCAST, DL, MVT::v2i64, Op->getOperand(2));
SDValue Lo = DAG.getNode(ISD::EXTRACT_VECTOR_ELT, DL, MVT::i64, Cast,
- DAG.getIntPtrConstant(0, DL));
+ DAG.getIntPtrConstant(0, DL));
SDValue Hi = DAG.getNode(ISD::EXTRACT_VECTOR_ELT, DL, MVT::i64, Cast,
- DAG.getIntPtrConstant(1, DL));
+ DAG.getIntPtrConstant(1, DL));
SmallVector<SDValue, 8> NewOps(Op->getNumOperands() + 1);
SmallVector<EVT, 8> ResultsType(Node->value_begin(), Node->value_end());
@@ -3130,11 +3130,12 @@ SDValue NVPTXTargetLowering::LowerCopyToReg_128(SDValue Op,
return DAG.getNode(ISD::CopyToReg, DL, ResultsType, NewOps);
}
-unsigned NVPTXTargetLowering::getNumRegisters(LLVMContext &Context, EVT VT,
- std::optional<MVT> RegisterVT = std::nullopt) const {
- if(VT == MVT::i128 && RegisterVT == MVT::i128)
- return 1;
- return TargetLoweringBase::getNumRegisters(Context, VT, RegisterVT);
+unsigned NVPTXTargetLowering::getNumRegisters(
+ LLVMContext &Context, EVT VT,
+ std::optional<MVT> RegisterVT = std::nullopt) const {
+ if (VT == MVT::i128 && RegisterVT == MVT::i128)
+ return 1;
+ return TargetLoweringBase::getNumRegisters(Context, VT, RegisterVT);
}
bool NVPTXTargetLowering::splitValueIntoRegisterParts(
@@ -5205,7 +5206,7 @@ NVPTXTargetLowering::getConstraintType(StringRef Constraint) const {
case 'l':
case 'f':
case 'd':
- case 'q':
+ case 'q':
case '0':
case 'N':
return C_RegisterClass;
@@ -6306,8 +6307,8 @@ static void ReplaceINTRINSIC_W_CHAIN(SDNode *N, SelectionDAG &DAG,
static void ReplaceCopyFromReg_128(SDNode *N, SelectionDAG &DAG,
SmallVectorImpl<SDValue> &Results) {
- // Change the CopyFromReg to output 2 64-bit results instead of a 128-bit result
- // so that it can pass the legalization
+ // Change the CopyFromReg to output 2 64-bit results instead of a 128-bit
+ // result so that it can pass the legalization
SDLoc DL(N);
SDValue Chain = N->getOperand(0);
SDValue Reg = N->getOperand(1);
diff --git a/llvm/test/CodeGen/NVPTX/inline-asm-b128-test1.ll b/llvm/test/CodeGen/NVPTX/inline-asm-b128-test1.ll
index dec0451c34ccc..8b5369d2804b0 100644
--- a/llvm/test/CodeGen/NVPTX/inline-asm-b128-test1.ll
+++ b/llvm/test/CodeGen/NVPTX/inline-asm-b128-test1.ll
@@ -3,22 +3,21 @@
target triple = "nvptx64-nvidia-cuda"
@value = internal addrspace(1) global i128 0, align 16
- at llvm.used = appending global [6 x ptr] [ptr @_Z7kernel1v, ptr @_Z7kernel2Pn, ptr @_Z7kernel3Pb, ptr @_Z7kernel4v, ptr @_Z7kernel5Pn, ptr addrspacecast (ptr addrspace(1) @value to ptr)], section "llvm.metadata"
; Function Attrs: alwaysinline convergent mustprogress willreturn
-define void @_Z7kernel1v() #0 {
+define void @_Z7kernel1v() {
; CHECK-LABEL: _Z7kernel1v
; CHECK: mov.u64 [[REG_HI:%rd[0-9]+]], 0;
; CHECK: mov.u64 [[REG_LO:%rd[0-9]+]], 42;
; CHECK: mov.b128 [[REG_128:%rq[0-9]+]], {[[REG_LO]], [[REG_HI]]};
; CHECK: { st.b128 [{{%rd[0-9]+}}], [[REG_128]]; }
- tail call void asm sideeffect "{ st.b128 [$0], $1; }", "l,q"(ptr nonnull addrspacecast (ptr addrspace(1) @value to ptr), i128 42) #3
+ tail call void asm sideeffect "{ st.b128 [$0], $1; }", "l,q"(ptr nonnull addrspacecast (ptr addrspace(1) @value to ptr), i128 42)
ret void
}
; Function Attrs: alwaysinline convergent mustprogress willreturn
-define void @_Z7kernel2Pn(ptr nocapture readonly %data) #0 {
+define void @_Z7kernel2Pn(ptr nocapture readonly %data) {
; CHECK-LABEL: _Z7kernel2Pn
; CHECK: ld.global.u64 [[REG_HI:%rd[0-9]+]], [[[REG_Addr:%r[0-9]+]]+8];
; CHECK: ld.global.u64 [[REG_LO:%rd[0-9]+]], [[[REG_Addr]]];
@@ -28,12 +27,12 @@ define void @_Z7kernel2Pn(ptr nocapture readonly %data) #0 {
%1 = addrspacecast ptr %data to ptr addrspace(1)
%2 = load <2 x i64>, ptr addrspace(1) %1, align 16
%3 = bitcast <2 x i64> %2 to i128
- tail call void asm sideeffect "{ st.b128 [$0], $1; }", "l,q"(ptr nonnull addrspacecast (ptr addrspace(1) @value to ptr), i128 %3) #3
+ tail call void asm sideeffect "{ st.b128 [$0], $1; }", "l,q"(ptr nonnull addrspacecast (ptr addrspace(1) @value to ptr), i128 %3)
ret void
}
; Function Attrs: alwaysinline convergent mustprogress willreturn
-define void @_Z7kernel3Pb(ptr nocapture readonly %flag) #0 {
+define void @_Z7kernel3Pb(ptr nocapture readonly %flag) {
; CHECK-LABEL: _Z7kernel3Pb
; CHECK: selp.b64 [[REG_LO:%rd[0-9]+]], 24, 42, {{%p[0-9]+}};
; CHECK: mov.u64 [[REG_HI:%rd[0-9]+]], 0;
@@ -44,17 +43,17 @@ define void @_Z7kernel3Pb(ptr nocapture readonly %flag) #0 {
%tmp1 = load i8, ptr addrspace(1) %1, align 1
%tobool.not = icmp eq i8 %tmp1, 0
%. = select i1 %tobool.not, i128 24, i128 42
- tail call void asm sideeffect "{ st.b128 [$0], $1; }", "l,q"(ptr nonnull addrspacecast (ptr addrspace(1) @value to ptr), i128 %.) #3
+ tail call void asm sideeffect "{ st.b128 [$0], $1; }", "l,q"(ptr nonnull addrspacecast (ptr addrspace(1) @value to ptr), i128 %.)
ret void
}
; Function Attrs: alwaysinline mustprogress willreturn memory(write, argmem: none, inaccessiblemem: none)
-define void @_Z7kernel4v() #1 {
+define void @_Z7kernel4v() {
; CHECK-LABEL: _Z7kernel4v
; CHECK-O3: { mov.b128 [[REG_128:%rq[0-9]+]], 41; }
; CHECK-O3: mov.b128 {%rd{{[0-9]+}}, %rd{{[0-9]+}}}, [[REG_128]];
- %1 = tail call i128 asm "{ mov.b128 $0, 41; }", "=q"() #4
+ %1 = tail call i128 asm "{ mov.b128 $0, 41; }", "=q"()
%add = add nsw i128 %1, 1
%2 = bitcast i128 %add to <2 x i64>
store <2 x i64> %2, ptr addrspace(1) @value, align 16
@@ -62,7 +61,7 @@ define void @_Z7kernel4v() #1 {
}
; Function Attrs: alwaysinline mustprogress willreturn memory(write, argmem: read, inaccessiblemem: none)
-define void @_Z7kernel5Pn(ptr nocapture readonly %data) #2 {
+define void @_Z7kernel5Pn(ptr nocapture readonly %data) {
; CHECK-LABEL: _Z7kernel5Pn
; CHECK-O3: ld.global.v2.u64 {[[REG_LO_IN:%rd[0-9]+]], [[REG_HI_IN:%rd[0-9]+]]}, [{{%rd[0-9]+}}];
; CHECK-O3: mov.b128 [[REG_128_IN:%rq[0-9]+]], {[[REG_LO_IN]], [[REG_HI_IN]]};
@@ -72,20 +71,13 @@ define void @_Z7kernel5Pn(ptr nocapture readonly %data) #2 {
%1 = addrspacecast ptr %data to ptr addrspace(1)
%2 = load <2 x i64>, ptr addrspace(1) %1, align 16
%3 = bitcast <2 x i64> %2 to i128
- %4 = tail call i128 asm "{ mov.b128 $0, $1; }", "=q,q"(i128 %3) #4
+ %4 = tail call i128 asm "{ mov.b128 $0, $1; }", "=q,q"(i128 %3)
%add = add nsw i128 %4, 1
%5 = bitcast i128 %add to <2 x i64>
store <2 x i64> %5, ptr addrspace(1) @value, align 16
ret void
}
-attributes #0 = { alwaysinline convergent mustprogress willreturn "nvvm.annotations_transplanted" "nvvm.kernel" "nvvm.restrict_processed" "target-cpu"="sm_89" }
-attributes #1 = { alwaysinline mustprogress willreturn memory(write, argmem: none, inaccessiblemem: none) "nvvm.annotations_transplanted" "nvvm.kernel" "nvvm.restrict_processed" "target-cpu"="sm_89" }
-attributes #2 = { alwaysinline mustprogress willreturn memory(write, argmem: read, inaccessiblemem: none) "nvvm.annotations_transplanted" "nvvm.kernel" "nvvm.restrict_processed" "target-cpu"="sm_89" }
-attributes #3 = { convergent nounwind }
-attributes #4 = { nounwind }
-
-
!nvvmir.version = !{!0, !1, !0, !1, !1, !0, !0, !0, !1}
!0 = !{i32 2, i32 0, i32 3, i32 1}
diff --git a/llvm/test/CodeGen/NVPTX/inline-asm-b128-test2.ll b/llvm/test/CodeGen/NVPTX/inline-asm-b128-test2.ll
index 337479a06c3f0..94b641e8faf05 100644
--- a/llvm/test/CodeGen/NVPTX/inline-asm-b128-test2.ll
+++ b/llvm/test/CodeGen/NVPTX/inline-asm-b128-test2.ll
@@ -10,48 +10,40 @@ target triple = "nvptx64-nvidia-cuda"
@v_u128_zero = internal addrspace(1) global i128 0, align 16
@v_i128_max = internal addrspace(1) global i128 0, align 16
@v_i128_min = internal addrspace(1) global i128 0, align 16
- at v64 = internal addrspace(1) global i64* null, align 8
- at llvm.used = appending global [10 x i8*] [i8* bitcast (i128* addrspacecast (i128 addrspace(1)* @u128_max to i128*) to i8*), i8* bitcast (i128* addrspacecast (i128 addrspace(1)* @u128_zero to i128*) to i8*), i8* bitcast (i128* addrspacecast (i128 addrspace(1)* @i128_max to i128*) to i8*), i8* bitcast (i128* addrspacecast (i128 addrspace(1)* @i128_min to i128*) to i8*), i8* bitcast (i128* addrspacecast (i128 addrspace(1)* @v_u128_max to i128*) to i8*), i8* bitcast (i128* addrspacecast (i128 addrspace(1)* @v_u128_zero to i128*) to i8*), i8* bitcast (i128* addrspacecast (i128 addrspace(1)* @v_i128_max to i128*) to i8*), i8* bitcast (i128* addrspacecast (i128 addrspace(1)* @v_i128_min to i128*) to i8*), i8* bitcast (i64** addrspacecast (i64* addrspace(1)* @v64 to i64**) to i8*), i8* bitcast (void ()* @_Z6kernelv to i8*)], section "llvm.metadata"
+ at v64 = internal addrspace(1) global ptr null, align 8
-; Function Attrs: alwaysinline
-define void @_Z6kernelv() #0 {
- ; CHECK-LABLE: _Z6kernelv
- ; CHECK: mov.u64 [[U64_MAX:%rd[0-9]+]], -1;
- ; CHECK: mov.b128 [[U128_MAX:%rq[0-9]+]], {[[U64_MAX]], [[U64_MAX]]};
- ; CHECK: mov.u64 [[I128_MAX_HI:%rd[0-9]+]], 9223372036854775807;
- ; CHECK: mov.b128 [[I128_MAX:%rq[0-9]+]], {[[U64_MAX]], [[I128_MAX_HI]]};
- ; CHECK: mov.u64 [[I128_MIN_HI:%rd[0-9]+]], -9223372036854775808;
- ; CHECK: mov.u64 [[ZERO:%rd[0-9]+]], 0;
- ; CHECK: mov.b128 [[I128_MIN:%rq[0-9]+]], {[[ZERO]], [[I128_MIN_HI]]};
- ; CHECK: mov.b128 [[U128_ZERO:%rq[0-9]+]], {[[ZERO]], [[ZERO]]};
-
- %tmp = load i64*, i64** addrspacecast (i64* addrspace(1)* @v64 to i64**), align 8
- %add.ptr = getelementptr inbounds i64, i64* %tmp, i32 0
- %tmp1 = load i64*, i64** addrspacecast (i64* addrspace(1)* @v64 to i64**), align 8
- %add.ptr2 = getelementptr inbounds i64, i64* %tmp1, i32 1
- call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}\0A\09", "q,l,l,l"(i128 -1, i64* %add.ptr, i64* %add.ptr2, i128* addrspacecast (i128 addrspace(1)* @v_u128_max to i128*)) #1
- %tmp3 = load i64*, i64** addrspacecast (i64* addrspace(1)* @v64 to i64**), align 8
- %add.ptr4 = getelementptr inbounds i64, i64* %tmp3, i32 2
- %tmp5 = load i64*, i64** addrspacecast (i64* addrspace(1)* @v64 to i64**), align 8
- %add.ptr6 = getelementptr inbounds i64, i64* %tmp5, i32 3
- call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}\0A\09", "q,l,l,l"(i128 170141183460469231731687303715884105727, i64* %add.ptr4, i64* %add.ptr6, i128* addrspacecast (i128 addrspace(1)* @v_i128_max to i128*)) #1
- %tmp7 = load i64*, i64** addrspacecast (i64* addrspace(1)* @v64 to i64**), align 8
- %add.ptr8 = getelementptr inbounds i64, i64* %tmp7, i32 4
- %tmp9 = load i64*, i64** addrspacecast (i64* addrspace(1)* @v64 to i64**), align 8
- %add.ptr10 = getelementptr inbounds i64, i64* %tmp9, i32 5
- call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}\0A\09", "q,l,l,l"(i128 -170141183460469231731687303715884105728, i64* %add.ptr8, i64* %add.ptr10, i128* addrspacecast (i128 addrspace(1)* @v_i128_min to i128*)) #1
- %tmp11 = load i64*, i64** addrspacecast (i64* addrspace(1)* @v64 to i64**), align 8
- %add.ptr12 = getelementptr inbounds i64, i64* %tmp11, i32 6
- %tmp13 = load i64*, i64** addrspacecast (i64* addrspace(1)* @v64 to i64**), align 8
- %add.ptr14 = getelementptr inbounds i64, i64* %tmp13, i32 7
- call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}\0A\09", "q,l,l,l"(i128 0, i64* %add.ptr12, i64* %add.ptr14, i128* addrspacecast (i128 addrspace(1)* @v_u128_zero to i128*)) #1
+; Function Attrs: alwaysinline convergent mustprogress willreturn
+define void @_Z6kernelv() {
+ ; CHECK-LABEL: _Z6kernelv
+ ; CHECK: mov.u64 [[U64_MAX:%rd[0-9]+]], -1;
+ ; CHECK: mov.b128 [[U128_MAX:%rq[0-9]+]], {[[U64_MAX]], [[U64_MAX]]};
+ ; CHECK: mov.u64 [[I64_MAX:%rd[0-9]+]], 9223372036854775807;
+ ; CHECK: mov.b128 [[I128_MAX:%rq[0-9]+]], {[[U64_MAX]], [[I64_MAX]]}
+ ; CHECK: mov.u64 [[I64_MIN:%rd[0-9]+]], -9223372036854775808;
+ ; CHECK: mov.u64 [[U64_ZERO:%rd[0-9]+]], 0;
+ ; CHECK: mov.b128 [[I128_MIN:%rq[0-9]+]], {[[U64_ZERO]], [[I64_MIN]]}
+ ; CHECK: mov.b128 [[U128_ZERO:%rq[0-9]+]], {[[U64_ZERO]], [[U64_ZERO]]}
+
+ %tmp = load ptr, ptr addrspace(1) @v64, align 8
+ %add.ptr2 = getelementptr inbounds i64, ptr %tmp, i64 1
+ tail call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}\0A\09", "q,l,l,l"(i128 -1, ptr %tmp, ptr nonnull %add.ptr2, ptr nonnull addrspacecast (ptr addrspace(1) @v_u128_max to ptr))
+ %tmp3 = load ptr, ptr addrspace(1) @v64, align 8
+ %add.ptr4 = getelementptr inbounds i64, ptr %tmp3, i64 2
+ %add.ptr6 = getelementptr inbounds i64, ptr %tmp3, i64 3
+ tail call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}\0A\09", "q,l,l,l"(i128 170141183460469231731687303715884105727, ptr nonnull %add.ptr4, ptr nonnull %add.ptr6, ptr nonnull addrspacecast (ptr addrspace(1) @v_i128_max to ptr))
+ %tmp7 = load ptr, ptr addrspace(1) @v64, align 8
+ %add.ptr8 = getelementptr inbounds i64, ptr %tmp7, i64 4
+ %add.ptr10 = getelementptr inbounds i64, ptr %tmp7, i64 5
+ tail call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}\0A\09", "q,l,l,l"(i128 -170141183460469231731687303715884105728, ptr nonnull %add.ptr8, ptr nonnull %add.ptr10, ptr nonnull addrspacecast (ptr addrspace(1) @v_i128_min to ptr))
+ %tmp11 = load ptr, ptr addrspace(1) @v64, align 8
+ %add.ptr12 = getelementptr inbounds i64, ptr %tmp11, i64 6
+ %add.ptr14 = getelementptr inbounds i64, ptr %tmp11, i64 7
+ tail call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}\0A\09", "q,l,l,l"(i128 0, ptr nonnull %add.ptr12, ptr nonnull %add.ptr14, ptr nonnull addrspacecast (ptr addrspace(1) @v_u128_zero to ptr))
ret void
}
-attributes #0 = { alwaysinline "nvvm.annotations_transplanted" "nvvm.kernel" }
-attributes #1 = { nounwind }
-!nvvmir.version = !{!0, !1, !0, !1, !1, !0, !0, !0, !1}
+!nvvmir.version = !{!2, !3, !2, !3, !3, !2, !2, !2, !3}
-!0 = !{i32 2, i32 0, i32 3, i32 1}
-!1 = !{i32 2, i32 0}
+!2 = !{i32 2, i32 0, i32 3, i32 1}
+!3 = !{i32 2, i32 0}
diff --git a/llvm/test/CodeGen/NVPTX/inline-asm-b128-test3.ll b/llvm/test/CodeGen/NVPTX/inline-asm-b128-test3.ll
index 4f077ec5383c9..9d7a25ca4d467 100644
--- a/llvm/test/CodeGen/NVPTX/inline-asm-b128-test3.ll
+++ b/llvm/test/CodeGen/NVPTX/inline-asm-b128-test3.ll
@@ -10,7 +10,7 @@ target triple = "nvptx64-nvidia-cuda"
@llvm.used = appending global [6 x ptr] [ptr @_Z6kernelv, ptr addrspacecast (ptr addrspace(1) @size to ptr), ptr addrspacecast (ptr addrspace(1) @value to ptr), ptr addrspacecast (ptr addrspace(1) @x to ptr), ptr addrspacecast (ptr addrspace(1) @y to ptr), ptr addrspacecast (ptr addrspace(1) @z to ptr)], section "llvm.metadata"
; Function Attrs: alwaysinline mustprogress willreturn memory(readwrite, argmem: none, inaccessiblemem: none)
-define void @_Z6kernelv() #0 {
+define void @_Z6kernelv() {
; CHECK-LABEL: _Z6kernelv
; CHECK: mov.b128 [[X:%rq[0-9]+]], {{{%rd[0-9]+}}, {{%rd[0-9]+}}};
; CHECK: mov.b128 [[Y:%rq[0-9]+]], {{{%rd[0-9]+}}, {{%rd[0-9]+}}};
@@ -62,43 +62,43 @@ define void @_Z6kernelv() #0 {
%5 = phi i128 [ %x.promoted5, %.lr.ph.preheader.new ], [ %asmresult19.3, %.lr.ph ]
%i.04 = phi i64 [ 0, %.lr.ph.preheader.new ], [ %inc.3, %.lr.ph ]
%niter = phi i64 [ 0, %.lr.ph.preheader.new ], [ %niter.next.3, %.lr.ph ]
- %6 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09add.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09add.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09add.cc.u64 lo, lo, 3;\0A\09add.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %i.04, i128 %5, i128 %4, i128 %3) #1
+ %6 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09add.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09add.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09add.cc.u64 lo, lo, 3;\0A\09add.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %i.04, i128 %5, i128 %4, i128 %3)
%asmresult = extractvalue { i128, i128, i128 } %6, 0
%asmresult7 = extractvalue { i128, i128, i128 } %6, 1
%asmresult8 = extractvalue { i128, i128, i128 } %6, 2
%add = add nsw i128 %asmresult, %asmresult7
%add12 = add nsw i128 %add, %asmresult8
%add14 = add nsw i128 %add12, %2
- %7 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09sub.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09sub.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09sub.cc.u64 lo, lo, 3;\0A\09sub.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %i.04, i128 %asmresult, i128 %asmresult7, i128 %asmresult8) #1
+ %7 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09sub.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09sub.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09sub.cc.u64 lo, lo, 3;\0A\09sub.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %i.04, i128 %asmresult, i128 %asmresult7, i128 %asmresult8)
%asmresult19 = extractvalue { i128, i128, i128 } %7, 0
%asmresult20 = extractvalue { i128, i128, i128 } %7, 1
%asmresult21 = extractvalue { i128, i128, i128 } %7, 2
%inc = add nuw nsw i64 %i.04, 1
- %8 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09add.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09add.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09add.cc.u64 lo, lo, 3;\0A\09add.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %inc, i128 %asmresult19, i128 %asmresult20, i128 %asmresult21) #1
+ %8 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09add.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09add.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09add.cc.u64 lo, lo, 3;\0A\09add.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %inc, i128 %asmresult19, i128 %asmresult20, i128 %asmresult21)
%asmresult.1 = extractvalue { i128, i128, i128 } %8, 0
%asmresult7.1 = extractvalue { i128, i128, i128 } %8, 1
%asmresult8.1 = extractvalue { i128, i128, i128 } %8, 2
%add.1 = add nsw i128 %asmresult.1, %asmresult7.1
%add12.1 = add nsw i128 %add.1, %asmresult8.1
%add14.1 = add nsw i128 %add12.1, %add14
- %9 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09sub.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09sub.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09sub.cc.u64 lo, lo, 3;\0A\09sub.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %inc, i128 %asmresult.1, i128 %asmresult7.1, i128 %asmresult8.1) #1
+ %9 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09sub.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09sub.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09sub.cc.u64 lo, lo, 3;\0A\09sub.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %inc, i128 %asmresult.1, i128 %asmresult7.1, i128 %asmresult8.1)
%asmresult19.1 = extractvalue { i128, i128, i128 } %9, 0
%asmresult20.1 = extractvalue { i128, i128, i128 } %9, 1
%asmresult21.1 = extractvalue { i128, i128, i128 } %9, 2
%inc.1 = add nuw nsw i64 %i.04, 2
- %10 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09add.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09add.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09add.cc.u64 lo, lo, 3;\0A\09add.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %inc.1, i128 %asmresult19.1, i128 %asmresult20.1, i128 %asmresult21.1) #1
+ %10 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09add.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09add.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09add.cc.u64 lo, lo, 3;\0A\09add.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %inc.1, i128 %asmresult19.1, i128 %asmresult20.1, i128 %asmresult21.1)
%asmresult.2 = extractvalue { i128, i128, i128 } %10, 0
%asmresult7.2 = extractvalue { i128, i128, i128 } %10, 1
%asmresult8.2 = extractvalue { i128, i128, i128 } %10, 2
%add.2 = add nsw i128 %asmresult.2, %asmresult7.2
%add12.2 = add nsw i128 %add.2, %asmresult8.2
%add14.2 = add nsw i128 %add12.2, %add14.1
- %11 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09sub.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09sub.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09sub.cc.u64 lo, lo, 3;\0A\09sub.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %inc.1, i128 %asmresult.2, i128 %asmresult7.2, i128 %asmresult8.2) #1
+ %11 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09sub.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09sub.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09sub.cc.u64 lo, lo, 3;\0A\09sub.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %inc.1, i128 %asmresult.2, i128 %asmresult7.2, i128 %asmresult8.2)
%asmresult19.2 = extractvalue { i128, i128, i128 } %11, 0
%asmresult20.2 = extractvalue { i128, i128, i128 } %11, 1
%asmresult21.2 = extractvalue { i128, i128, i128 } %11, 2
%inc.2 = add nuw nsw i64 %i.04, 3
- %12 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09add.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09add.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09add.cc.u64 lo, lo, 3;\0A\09add.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %inc.2, i128 %asmresult19.2, i128 %asmresult20.2, i128 %asmresult21.2) #1
+ %12 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09add.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09add.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09add.cc.u64 lo, lo, 3;\0A\09add.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %inc.2, i128 %asmresult19.2, i128 %asmresult20.2, i128 %asmresult21.2)
%asmresult.3 = extractvalue { i128, i128, i128 } %12, 0
%asmresult7.3 = extractvalue { i128, i128, i128 } %12, 1
%asmresult8.3 = extractvalue { i128, i128, i128 } %12, 2
@@ -107,7 +107,7 @@ define void @_Z6kernelv() #0 {
%add14.3 = add nsw i128 %add12.3, %add14.2
%13 = bitcast i128 %add14.3 to <2 x i64>
store <2 x i64> %13, ptr addrspace(1) @value, align 16
- %14 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09sub.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09sub.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09sub.cc.u64 lo, lo, 3;\0A\09sub.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %inc.2, i128 %asmresult.3, i128 %asmresult7.3, i128 %asmresult8.3) #1
+ %14 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09sub.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09sub.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09sub.cc.u64 lo, lo, 3;\0A\09sub.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %inc.2, i128 %asmresult.3, i128 %asmresult7.3, i128 %asmresult8.3)
%asmresult19.3 = extractvalue { i128, i128, i128 } %14, 0
%asmresult20.3 = extractvalue { i128, i128, i128 } %14, 1
%asmresult21.3 = extractvalue { i128, i128, i128 } %14, 2
@@ -138,7 +138,7 @@ define void @_Z6kernelv() #0 {
%21 = phi i128 [ %asmresult19.epil, %.lr.ph.epil ], [ %.unr11, %._crit_edge.loopexit.unr-lcssa ]
%i.04.epil = phi i64 [ %inc.epil, %.lr.ph.epil ], [ %i.04.unr, %._crit_edge.loopexit.unr-lcssa ]
%epil.iter = phi i64 [ %epil.iter.next, %.lr.ph.epil ], [ 0, %._crit_edge.loopexit.unr-lcssa ]
- %22 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09add.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09add.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09add.cc.u64 lo, lo, 3;\0A\09add.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %i.04.epil, i128 %21, i128 %20, i128 %19) #1
+ %22 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09add.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09add.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09add.cc.u64 lo, lo, 3;\0A\09add.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %i.04.epil, i128 %21, i128 %20, i128 %19)
%asmresult.epil = extractvalue { i128, i128, i128 } %22, 0
%asmresult7.epil = extractvalue { i128, i128, i128 } %22, 1
%asmresult8.epil = extractvalue { i128, i128, i128 } %22, 2
@@ -147,7 +147,7 @@ define void @_Z6kernelv() #0 {
%add14.epil = add nsw i128 %add12.epil, %18
%23 = bitcast i128 %add14.epil to <2 x i64>
store <2 x i64> %23, ptr addrspace(1) @value, align 16
- %24 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09sub.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09sub.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09sub.cc.u64 lo, lo, 3;\0A\09sub.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %i.04.epil, i128 %asmresult.epil, i128 %asmresult7.epil, i128 %asmresult8.epil) #1
+ %24 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09sub.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09sub.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09sub.cc.u64 lo, lo, 3;\0A\09sub.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %i.04.epil, i128 %asmresult.epil, i128 %asmresult7.epil, i128 %asmresult8.epil)
%asmresult19.epil = extractvalue { i128, i128, i128 } %24, 0
%asmresult20.epil = extractvalue { i128, i128, i128 } %24, 1
%asmresult21.epil = extractvalue { i128, i128, i128 } %24, 2
@@ -166,8 +166,6 @@ define void @_Z6kernelv() #0 {
ret void
}
-attributes #0 = { alwaysinline mustprogress willreturn memory(readwrite, argmem: none, inaccessiblemem: none) "nvvm.annotations_transplanted" "nvvm.kernel" "nvvm.restrict_processed" "target-cpu"="sm_89" }
-attributes #1 = { nounwind }
!nvvmir.version = !{!0, !1, !0, !1, !1, !0, !0, !0, !1}
>From 7274c66310656804df9d4f698fa3013c28b5db36 Mon Sep 17 00:00:00 2001
From: chengjunp <chengjunp at nvidia.com>
Date: Wed, 26 Jun 2024 00:04:10 +0000
Subject: [PATCH 3/6] Update kernel names in tests & Update one test for inline
asm in loops
---
.../CodeGen/NVPTX/inline-asm-b128-test1.ll | 36 ++--
.../CodeGen/NVPTX/inline-asm-b128-test2.ll | 7 +-
.../CodeGen/NVPTX/inline-asm-b128-test3.ll | 177 +++---------------
3 files changed, 45 insertions(+), 175 deletions(-)
diff --git a/llvm/test/CodeGen/NVPTX/inline-asm-b128-test1.ll b/llvm/test/CodeGen/NVPTX/inline-asm-b128-test1.ll
index 8b5369d2804b0..8a256d50d6050 100644
--- a/llvm/test/CodeGen/NVPTX/inline-asm-b128-test1.ll
+++ b/llvm/test/CodeGen/NVPTX/inline-asm-b128-test1.ll
@@ -1,12 +1,13 @@
-; RUN: llc < %s -march=nvptx -mcpu=sm_70 -o - 2>&1 | FileCheck %s
+; RUN: llc < %s -march=nvptx -mcpu=sm_70 -mattr=+ptx83 | FileCheck %s
+; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_70 -mattr=+ptx83 | %ptxas-verify -arch=sm_70 %}
target triple = "nvptx64-nvidia-cuda"
@value = internal addrspace(1) global i128 0, align 16
; Function Attrs: alwaysinline convergent mustprogress willreturn
-define void @_Z7kernel1v() {
- ; CHECK-LABEL: _Z7kernel1v
+define void @test_b128_input_from_const() {
+ ; CHECK-LABEL: test_b128_input_from_const
; CHECK: mov.u64 [[REG_HI:%rd[0-9]+]], 0;
; CHECK: mov.u64 [[REG_LO:%rd[0-9]+]], 42;
; CHECK: mov.b128 [[REG_128:%rq[0-9]+]], {[[REG_LO]], [[REG_HI]]};
@@ -17,8 +18,8 @@ define void @_Z7kernel1v() {
}
; Function Attrs: alwaysinline convergent mustprogress willreturn
-define void @_Z7kernel2Pn(ptr nocapture readonly %data) {
- ; CHECK-LABEL: _Z7kernel2Pn
+define void @test_b128_input_from_load(ptr nocapture readonly %data) {
+ ; CHECK-LABEL: test_b128_input_from_load
; CHECK: ld.global.u64 [[REG_HI:%rd[0-9]+]], [[[REG_Addr:%r[0-9]+]]+8];
; CHECK: ld.global.u64 [[REG_LO:%rd[0-9]+]], [[[REG_Addr]]];
; CHECK: mov.b128 [[REG_128:%rq[0-9]+]], {[[REG_LO]], [[REG_HI]]};
@@ -32,8 +33,8 @@ define void @_Z7kernel2Pn(ptr nocapture readonly %data) {
}
; Function Attrs: alwaysinline convergent mustprogress willreturn
-define void @_Z7kernel3Pb(ptr nocapture readonly %flag) {
- ; CHECK-LABEL: _Z7kernel3Pb
+define void @test_b128_input_from_select(ptr nocapture readonly %flag) {
+ ; CHECK-LABEL: test_b128_input_from_select
; CHECK: selp.b64 [[REG_LO:%rd[0-9]+]], 24, 42, {{%p[0-9]+}};
; CHECK: mov.u64 [[REG_HI:%rd[0-9]+]], 0;
; CHECK: mov.b128 [[REG_128:%rq[0-9]+]], {[[REG_LO]], [[REG_HI]]};
@@ -48,10 +49,10 @@ define void @_Z7kernel3Pb(ptr nocapture readonly %flag) {
}
; Function Attrs: alwaysinline mustprogress willreturn memory(write, argmem: none, inaccessiblemem: none)
-define void @_Z7kernel4v() {
- ; CHECK-LABEL: _Z7kernel4v
- ; CHECK-O3: { mov.b128 [[REG_128:%rq[0-9]+]], 41; }
- ; CHECK-O3: mov.b128 {%rd{{[0-9]+}}, %rd{{[0-9]+}}}, [[REG_128]];
+define void @test_store_b128_output() {
+ ; CHECK-LABEL: test_store_b128_output
+ ; CHECK: { mov.b128 [[REG_128:%rq[0-9]+]], 41; }
+ ; CHECK: mov.b128 {%rd{{[0-9]+}}, %rd{{[0-9]+}}}, [[REG_128]];
%1 = tail call i128 asm "{ mov.b128 $0, 41; }", "=q"()
%add = add nsw i128 %1, 1
@@ -61,12 +62,13 @@ define void @_Z7kernel4v() {
}
; Function Attrs: alwaysinline mustprogress willreturn memory(write, argmem: read, inaccessiblemem: none)
-define void @_Z7kernel5Pn(ptr nocapture readonly %data) {
- ; CHECK-LABEL: _Z7kernel5Pn
- ; CHECK-O3: ld.global.v2.u64 {[[REG_LO_IN:%rd[0-9]+]], [[REG_HI_IN:%rd[0-9]+]]}, [{{%rd[0-9]+}}];
- ; CHECK-O3: mov.b128 [[REG_128_IN:%rq[0-9]+]], {[[REG_LO_IN]], [[REG_HI_IN]]};
- ; CHECK-O3: { mov.b128 [[REG_128_OUT:%rq[0-9]+]], [[REG_128_IN]]; }
- ; CHECK-O3: mov.b128 {%rd{{[0-9]+}}, %rd{{[0-9]+}}}, [[REG_128_OUT]];
+define void @test_use_of_b128_output(ptr nocapture readonly %data) {
+ ; CHECK-LABEL: test_use_of_b128_output
+ ; CHECK: ld.global.u64 [[REG_HI:%rd[0-9]+]], [[[REG_Addr:%r[0-9]+]]+8];
+ ; CHECK: ld.global.u64 [[REG_LO:%rd[0-9]+]], [[[REG_Addr]]];
+ ; CHECK: mov.b128 [[REG_128_IN:%rq[0-9]+]], {[[REG_LO]], [[REG_HI]]};
+ ; CHECK: { mov.b128 [[REG_128_OUT:%rq[0-9]+]], [[REG_128_IN]]; }
+ ; CHECK: mov.b128 {%rd{{[0-9]+}}, %rd{{[0-9]+}}}, [[REG_128_OUT]];
%1 = addrspacecast ptr %data to ptr addrspace(1)
%2 = load <2 x i64>, ptr addrspace(1) %1, align 16
diff --git a/llvm/test/CodeGen/NVPTX/inline-asm-b128-test2.ll b/llvm/test/CodeGen/NVPTX/inline-asm-b128-test2.ll
index 94b641e8faf05..09b648d036a4c 100644
--- a/llvm/test/CodeGen/NVPTX/inline-asm-b128-test2.ll
+++ b/llvm/test/CodeGen/NVPTX/inline-asm-b128-test2.ll
@@ -1,4 +1,5 @@
-; RUN: llc < %s -march=nvptx -mcpu=sm_70 -o - 2>&1 | FileCheck %s
+; RUN: llc < %s -march=nvptx -mcpu=sm_70 -mattr=+ptx83 | FileCheck %s
+; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_70 -mattr=+ptx83 | %ptxas-verify -arch=sm_70 %}
target triple = "nvptx64-nvidia-cuda"
@@ -13,8 +14,8 @@ target triple = "nvptx64-nvidia-cuda"
@v64 = internal addrspace(1) global ptr null, align 8
; Function Attrs: alwaysinline convergent mustprogress willreturn
-define void @_Z6kernelv() {
- ; CHECK-LABEL: _Z6kernelv
+define void @test_corner_values() {
+ ; CHECK-LABEL: test_corner_values
; CHECK: mov.u64 [[U64_MAX:%rd[0-9]+]], -1;
; CHECK: mov.b128 [[U128_MAX:%rq[0-9]+]], {[[U64_MAX]], [[U64_MAX]]};
; CHECK: mov.u64 [[I64_MAX:%rd[0-9]+]], 9223372036854775807;
diff --git a/llvm/test/CodeGen/NVPTX/inline-asm-b128-test3.ll b/llvm/test/CodeGen/NVPTX/inline-asm-b128-test3.ll
index 9d7a25ca4d467..e187aa4370858 100644
--- a/llvm/test/CodeGen/NVPTX/inline-asm-b128-test3.ll
+++ b/llvm/test/CodeGen/NVPTX/inline-asm-b128-test3.ll
@@ -1,177 +1,44 @@
-; RUN: llc < %s -march=nvptx -mcpu=sm_70 -o - 2>&1 | FileCheck %s
+; RUN: llc < %s -march=nvptx -mcpu=sm_70 -mattr=+ptx83 | FileCheck %s
+; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_70 -mattr=+ptx83 | %ptxas-verify -arch=sm_70 %}
target triple = "nvptx64-nvidia-cuda"
@size = internal addrspace(1) global i32 0, align 4
- at value = internal addrspace(1) global i128 0, align 16
@x = internal addrspace(1) global i128 0, align 16
- at y = internal addrspace(1) global i128 0, align 16
- at z = internal addrspace(1) global i128 0, align 16
- at llvm.used = appending global [6 x ptr] [ptr @_Z6kernelv, ptr addrspacecast (ptr addrspace(1) @size to ptr), ptr addrspacecast (ptr addrspace(1) @value to ptr), ptr addrspacecast (ptr addrspace(1) @x to ptr), ptr addrspacecast (ptr addrspace(1) @y to ptr), ptr addrspacecast (ptr addrspace(1) @z to ptr)], section "llvm.metadata"
-; Function Attrs: alwaysinline mustprogress willreturn memory(readwrite, argmem: none, inaccessiblemem: none)
-define void @_Z6kernelv() {
- ; CHECK-LABEL: _Z6kernelv
- ; CHECK: mov.b128 [[X:%rq[0-9]+]], {{{%rd[0-9]+}}, {{%rd[0-9]+}}};
- ; CHECK: mov.b128 [[Y:%rq[0-9]+]], {{{%rd[0-9]+}}, {{%rd[0-9]+}}};
- ; CHECK: mov.b128 [[Z:%rq[0-9]+]], {{{%rd[0-9]+}}, {{%rd[0-9]+}}};
- ; CHECK: mov.b128 {lo, hi}, [[X]];
- ; CHECK: mov.b128 [[X]], {lo, hi};
- ; CHECK: mov.b128 {lo, hi}, [[Y]];
- ; CHECK: mov.b128 [[Y]], {lo, hi};
- ; CHECK: mov.b128 {lo, hi}, [[Z]];
- ; CHECK: mov.b128 [[Z]], {lo, hi};
- ; CHECK: mov.b128 {[[X_LO:%rd[0-9]+]], [[X_HI:%rd[0-9]+]]}, [[X]];
- ; CHECK: mov.b128 {[[Y_LO:%rd[0-9]+]], [[Y_HI:%rd[0-9]+]]}, [[Y]];
- ; CHECK: mov.b128 {[[Z_LO:%rd[0-9]+]], [[Z_HI:%rd[0-9]+]]}, [[Z]];
- ; CHECK: mov.b128 [[X_NEW:%rq[0-9]+]], {[[X_LO]], [[X_HI]]};
- ; CHECK: mov.b128 [[Y_NEW:%rq[0-9]+]], {[[Y_LO]], [[Y_HI]]};
- ; CHECK: mov.b128 [[Z_NEW:%rq[0-9]+]], {[[Z_LO]], [[Z_HI]]};
- ; CHECK: mov.b128 {lo, hi}, [[X_NEW]];
- ; CHECK: mov.b128 [[X_NEW]], {lo, hi};
- ; CHECK: mov.b128 {lo, hi}, [[Y_NEW]];
- ; CHECK: mov.b128 [[Y_NEW]], {lo, hi};
- ; CHECK: mov.b128 {lo, hi}, [[Z_NEW]];
- ; CHECK: mov.b128 [[Z_NEW]], {lo, hi};
- ; CHECK: mov.b128 {{{%rd[0-9]+}}, {{%rd[0-9]+}}}, [[X_NEW]];
- ; CHECK: mov.b128 {{{%rd[0-9]+}}, {{%rd[0-9]+}}}, [[Y_NEW]];
- ; CHECK: mov.b128 {{{%rd[0-9]+}}, {{%rd[0-9]+}}}, [[Z_NEW]];
-
+define void @test_b128_in_loop() {
+ ; CHECK-LABEL: test_b128_in_loop
+ ; CHECK: ld.global.u64 [[REG_HI:%rd[0-9]+]], [x+8];
+ ; CHECK: ld.global.u64 [[REG_LO:%rd[0-9]+]], [x];
+ ; CHECK: mov.b128 [[REG_128:%rq[0-9]+]], {[[REG_LO]], [[REG_HI]]};
+ ; CHECK: mov.b128 {lo, hi}, [[REG_128]];
+ ; CHECK: add.cc.u64 lo, lo, {{%rd[0-9]+}};
+ ; CHECK: mov.b128 [[REG_128]], {lo, hi};
+
%tmp11 = load i32, ptr addrspace(1) @size, align 4
%cmp3.not = icmp eq i32 %tmp11, 0
br i1 %cmp3.not, label %._crit_edge, label %.lr.ph.preheader
.lr.ph.preheader: ; preds = %0
%x.promoted5 = load i128, ptr addrspace(1) @x, align 16
- %y.promoted6 = load i128, ptr addrspace(1) @y, align 16
- %z.promoted7 = load i128, ptr addrspace(1) @z, align 16
- %value.promoted8 = load i128, ptr addrspace(1) @value, align 16
%umax = sext i32 %tmp11 to i64
- %xtraiter = and i64 %umax, 3
- %1 = icmp ult i32 %tmp11, 4
- br i1 %1, label %._crit_edge.loopexit.unr-lcssa, label %.lr.ph.preheader.new
-
-.lr.ph.preheader.new: ; preds = %.lr.ph.preheader
- %unroll_iter = and i64 %umax, -4
br label %.lr.ph
-.lr.ph: ; preds = %.lr.ph, %.lr.ph.preheader.new
- %2 = phi i128 [ %value.promoted8, %.lr.ph.preheader.new ], [ %add14.3, %.lr.ph ]
- %3 = phi i128 [ %z.promoted7, %.lr.ph.preheader.new ], [ %asmresult21.3, %.lr.ph ]
- %4 = phi i128 [ %y.promoted6, %.lr.ph.preheader.new ], [ %asmresult20.3, %.lr.ph ]
- %5 = phi i128 [ %x.promoted5, %.lr.ph.preheader.new ], [ %asmresult19.3, %.lr.ph ]
- %i.04 = phi i64 [ 0, %.lr.ph.preheader.new ], [ %inc.3, %.lr.ph ]
- %niter = phi i64 [ 0, %.lr.ph.preheader.new ], [ %niter.next.3, %.lr.ph ]
- %6 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09add.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09add.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09add.cc.u64 lo, lo, 3;\0A\09add.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %i.04, i128 %5, i128 %4, i128 %3)
- %asmresult = extractvalue { i128, i128, i128 } %6, 0
- %asmresult7 = extractvalue { i128, i128, i128 } %6, 1
- %asmresult8 = extractvalue { i128, i128, i128 } %6, 2
- %add = add nsw i128 %asmresult, %asmresult7
- %add12 = add nsw i128 %add, %asmresult8
- %add14 = add nsw i128 %add12, %2
- %7 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09sub.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09sub.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09sub.cc.u64 lo, lo, 3;\0A\09sub.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %i.04, i128 %asmresult, i128 %asmresult7, i128 %asmresult8)
- %asmresult19 = extractvalue { i128, i128, i128 } %7, 0
- %asmresult20 = extractvalue { i128, i128, i128 } %7, 1
- %asmresult21 = extractvalue { i128, i128, i128 } %7, 2
- %inc = add nuw nsw i64 %i.04, 1
- %8 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09add.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09add.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09add.cc.u64 lo, lo, 3;\0A\09add.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %inc, i128 %asmresult19, i128 %asmresult20, i128 %asmresult21)
- %asmresult.1 = extractvalue { i128, i128, i128 } %8, 0
- %asmresult7.1 = extractvalue { i128, i128, i128 } %8, 1
- %asmresult8.1 = extractvalue { i128, i128, i128 } %8, 2
- %add.1 = add nsw i128 %asmresult.1, %asmresult7.1
- %add12.1 = add nsw i128 %add.1, %asmresult8.1
- %add14.1 = add nsw i128 %add12.1, %add14
- %9 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09sub.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09sub.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09sub.cc.u64 lo, lo, 3;\0A\09sub.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %inc, i128 %asmresult.1, i128 %asmresult7.1, i128 %asmresult8.1)
- %asmresult19.1 = extractvalue { i128, i128, i128 } %9, 0
- %asmresult20.1 = extractvalue { i128, i128, i128 } %9, 1
- %asmresult21.1 = extractvalue { i128, i128, i128 } %9, 2
- %inc.1 = add nuw nsw i64 %i.04, 2
- %10 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09add.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09add.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09add.cc.u64 lo, lo, 3;\0A\09add.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %inc.1, i128 %asmresult19.1, i128 %asmresult20.1, i128 %asmresult21.1)
- %asmresult.2 = extractvalue { i128, i128, i128 } %10, 0
- %asmresult7.2 = extractvalue { i128, i128, i128 } %10, 1
- %asmresult8.2 = extractvalue { i128, i128, i128 } %10, 2
- %add.2 = add nsw i128 %asmresult.2, %asmresult7.2
- %add12.2 = add nsw i128 %add.2, %asmresult8.2
- %add14.2 = add nsw i128 %add12.2, %add14.1
- %11 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09sub.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09sub.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09sub.cc.u64 lo, lo, 3;\0A\09sub.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %inc.1, i128 %asmresult.2, i128 %asmresult7.2, i128 %asmresult8.2)
- %asmresult19.2 = extractvalue { i128, i128, i128 } %11, 0
- %asmresult20.2 = extractvalue { i128, i128, i128 } %11, 1
- %asmresult21.2 = extractvalue { i128, i128, i128 } %11, 2
- %inc.2 = add nuw nsw i64 %i.04, 3
- %12 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09add.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09add.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09add.cc.u64 lo, lo, 3;\0A\09add.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %inc.2, i128 %asmresult19.2, i128 %asmresult20.2, i128 %asmresult21.2)
- %asmresult.3 = extractvalue { i128, i128, i128 } %12, 0
- %asmresult7.3 = extractvalue { i128, i128, i128 } %12, 1
- %asmresult8.3 = extractvalue { i128, i128, i128 } %12, 2
- %add.3 = add nsw i128 %asmresult.3, %asmresult7.3
- %add12.3 = add nsw i128 %add.3, %asmresult8.3
- %add14.3 = add nsw i128 %add12.3, %add14.2
- %13 = bitcast i128 %add14.3 to <2 x i64>
- store <2 x i64> %13, ptr addrspace(1) @value, align 16
- %14 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09sub.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09sub.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09sub.cc.u64 lo, lo, 3;\0A\09sub.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %inc.2, i128 %asmresult.3, i128 %asmresult7.3, i128 %asmresult8.3)
- %asmresult19.3 = extractvalue { i128, i128, i128 } %14, 0
- %asmresult20.3 = extractvalue { i128, i128, i128 } %14, 1
- %asmresult21.3 = extractvalue { i128, i128, i128 } %14, 2
- %15 = bitcast i128 %asmresult19.3 to <2 x i64>
- store <2 x i64> %15, ptr addrspace(1) @x, align 16
- %16 = bitcast i128 %asmresult20.3 to <2 x i64>
- store <2 x i64> %16, ptr addrspace(1) @y, align 16
- %17 = bitcast i128 %asmresult21.3 to <2 x i64>
- store <2 x i64> %17, ptr addrspace(1) @z, align 16
- %inc.3 = add nuw i64 %i.04, 4
- %niter.next.3 = add i64 %niter, 4
- %niter.ncmp.3.not = icmp eq i64 %niter.next.3, %unroll_iter
- br i1 %niter.ncmp.3.not, label %._crit_edge.loopexit.unr-lcssa, label %.lr.ph, !llvm.loop !2
-
-._crit_edge.loopexit.unr-lcssa: ; preds = %.lr.ph, %.lr.ph.preheader
- %.unr = phi i128 [ %value.promoted8, %.lr.ph.preheader ], [ %add14.3, %.lr.ph ]
- %.unr9 = phi i128 [ %z.promoted7, %.lr.ph.preheader ], [ %asmresult21.3, %.lr.ph ]
- %.unr10 = phi i128 [ %y.promoted6, %.lr.ph.preheader ], [ %asmresult20.3, %.lr.ph ]
- %.unr11 = phi i128 [ %x.promoted5, %.lr.ph.preheader ], [ %asmresult19.3, %.lr.ph ]
- %i.04.unr = phi i64 [ 0, %.lr.ph.preheader ], [ %inc.3, %.lr.ph ]
- %lcmp.mod.not = icmp eq i64 %xtraiter, 0
- br i1 %lcmp.mod.not, label %._crit_edge, label %.lr.ph.epil
-
-.lr.ph.epil: ; preds = %.lr.ph.epil, %._crit_edge.loopexit.unr-lcssa
- %18 = phi i128 [ %add14.epil, %.lr.ph.epil ], [ %.unr, %._crit_edge.loopexit.unr-lcssa ]
- %19 = phi i128 [ %asmresult21.epil, %.lr.ph.epil ], [ %.unr9, %._crit_edge.loopexit.unr-lcssa ]
- %20 = phi i128 [ %asmresult20.epil, %.lr.ph.epil ], [ %.unr10, %._crit_edge.loopexit.unr-lcssa ]
- %21 = phi i128 [ %asmresult19.epil, %.lr.ph.epil ], [ %.unr11, %._crit_edge.loopexit.unr-lcssa ]
- %i.04.epil = phi i64 [ %inc.epil, %.lr.ph.epil ], [ %i.04.unr, %._crit_edge.loopexit.unr-lcssa ]
- %epil.iter = phi i64 [ %epil.iter.next, %.lr.ph.epil ], [ 0, %._crit_edge.loopexit.unr-lcssa ]
- %22 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09add.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09add.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09add.cc.u64 lo, lo, 3;\0A\09add.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %i.04.epil, i128 %21, i128 %20, i128 %19)
- %asmresult.epil = extractvalue { i128, i128, i128 } %22, 0
- %asmresult7.epil = extractvalue { i128, i128, i128 } %22, 1
- %asmresult8.epil = extractvalue { i128, i128, i128 } %22, 2
- %add.epil = add nsw i128 %asmresult.epil, %asmresult7.epil
- %add12.epil = add nsw i128 %add.epil, %asmresult8.epil
- %add14.epil = add nsw i128 %add12.epil, %18
- %23 = bitcast i128 %add14.epil to <2 x i64>
- store <2 x i64> %23, ptr addrspace(1) @value, align 16
- %24 = tail call { i128, i128, i128 } asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09sub.cc.u64 lo, lo, 1;\0A\09mov.b128 $0, {lo, hi};\0A\09mov.b128 {lo, hi}, $1;\0A\09sub.cc.u64 hi, hi, 2;\0A\09mov.b128 $1, {lo, hi};\0A\09mov.b128 {lo, hi}, $2;\0A\09sub.cc.u64 lo, lo, 3;\0A\09sub.cc.u64 hi, hi, 3;\0A\09mov.b128 $2, {lo, hi};\0A\09}\0A\09", "=q,=q,=q,l,0,1,2"(i64 %i.04.epil, i128 %asmresult.epil, i128 %asmresult7.epil, i128 %asmresult8.epil)
- %asmresult19.epil = extractvalue { i128, i128, i128 } %24, 0
- %asmresult20.epil = extractvalue { i128, i128, i128 } %24, 1
- %asmresult21.epil = extractvalue { i128, i128, i128 } %24, 2
- %25 = bitcast i128 %asmresult19.epil to <2 x i64>
- store <2 x i64> %25, ptr addrspace(1) @x, align 16
- %26 = bitcast i128 %asmresult20.epil to <2 x i64>
- store <2 x i64> %26, ptr addrspace(1) @y, align 16
- %27 = bitcast i128 %asmresult21.epil to <2 x i64>
- store <2 x i64> %27, ptr addrspace(1) @z, align 16
- %inc.epil = add nuw i64 %i.04.epil, 1
- %epil.iter.next = add i64 %epil.iter, 1
- %epil.iter.cmp.not = icmp eq i64 %epil.iter.next, %xtraiter
- br i1 %epil.iter.cmp.not, label %._crit_edge, label %.lr.ph.epil, !llvm.loop !4
-
-._crit_edge: ; preds = %.lr.ph.epil, %._crit_edge.loopexit.unr-lcssa, %0
+.lr.ph: ; preds = %.lr.ph, %.lr.ph.preheader
+ %1 = phi i128 [ %2, %.lr.ph ], [ %x.promoted5, %.lr.ph.preheader ]
+ %i.04 = phi i64 [ %inc, %.lr.ph ], [ 0, %.lr.ph.preheader ]
+ %2 = tail call i128 asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09add.cc.u64 lo, lo, $1;\0A\09mov.b128 $0, {lo, hi};\0A\09} \0A\09", "=q,l,0"(i64 %i.04, i128 %1)
+ %3 = bitcast i128 %2 to <2 x i64>
+ store <2 x i64> %3, ptr addrspace(1) @x, align 16
+ %inc = add nuw i64 %i.04, 1
+ %exitcond.not = icmp eq i64 %inc, %umax
+ br i1 %exitcond.not, label %._crit_edge, label %.lr.ph
+
+._crit_edge: ; preds = %.lr.ph, %0
ret void
}
-
!nvvmir.version = !{!0, !1, !0, !1, !1, !0, !0, !0, !1}
!0 = !{i32 2, i32 0, i32 3, i32 1}
!1 = !{i32 2, i32 0}
-!2 = distinct !{!2, !3}
-!3 = !{!"llvm.loop.mustprogress"}
-!4 = distinct !{!4, !5}
-!5 = !{!"llvm.loop.unroll.disable"}
>From 2443c0bc80a94083c10672630f7d43d01b69afac Mon Sep 17 00:00:00 2001
From: chengjunp <chengjunp at nvidia.com>
Date: Wed, 26 Jun 2024 20:47:52 +0000
Subject: [PATCH 4/6] Update testcases with checks generated by
update_llc_test_checks.py
---
llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp | 5 +-
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 17 +--
.../CodeGen/NVPTX/inline-asm-b128-test1.ll | 127 +++++++++++++-----
.../CodeGen/NVPTX/inline-asm-b128-test2.ll | 106 +++++++++++++--
.../CodeGen/NVPTX/inline-asm-b128-test3.ll | 44 ++++--
5 files changed, 232 insertions(+), 67 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
index 9c0498560db21..11193c11ede3b 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
@@ -3830,14 +3830,13 @@ void NVPTXDAGToDAGISel::SelectV2I64toI128(SDNode *N) {
SDNode *Mov =
CurDAG->getMachineNode(NVPTX::V2I64toI128, DL, MVT::i128, {Lo, Hi});
- SmallVector<EVT, 8> ResultsType(N->value_begin(), N->value_end());
- SmallVector<SDValue, 8> NewOps(N->getNumOperands() - 1);
+ SmallVector<SDValue, 4> NewOps(N->getNumOperands() - 1);
NewOps[0] = N->getOperand(0);
NewOps[1] = Dst;
NewOps[2] = SDValue(Mov, 0);
if (N->getNumOperands() == 5)
NewOps[3] = N->getOperand(4);
- SDValue NewValue = CurDAG->getNode(ISD::CopyToReg, DL, ResultsType, NewOps);
+ SDValue NewValue = CurDAG->getNode(ISD::CopyToReg, DL, SmallVector<EVT>(N->values()), NewOps);
ReplaceNode(N, NewValue.getNode());
}
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 3cfac3baeb5c9..fd4293e56cfb6 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -3117,8 +3117,8 @@ SDValue NVPTXTargetLowering::LowerCopyToReg_128(SDValue Op,
SDValue Hi = DAG.getNode(ISD::EXTRACT_VECTOR_ELT, DL, MVT::i64, Cast,
DAG.getIntPtrConstant(1, DL));
- SmallVector<SDValue, 8> NewOps(Op->getNumOperands() + 1);
- SmallVector<EVT, 8> ResultsType(Node->value_begin(), Node->value_end());
+ SmallVector<SDValue, 5> NewOps(Op->getNumOperands() + 1);
+ SmallVector<EVT, 3> ResultsType(Node->value_begin(), Node->value_end());
NewOps[0] = Op->getOperand(0); // Chain
NewOps[1] = Op->getOperand(1); // Dst Reg
@@ -6316,16 +6316,9 @@ static void ReplaceCopyFromReg_128(SDNode *N, SelectionDAG &DAG,
assert(Reg.getValueType() == MVT::i128 &&
"Custom lowering for CopyFromReg with 128-bit reg only");
- SmallVector<EVT, 8> ResultsType(4);
- SmallVector<SDValue, 8> NewOps(3);
- ResultsType[0] = MVT::i64;
- ResultsType[1] = MVT::i64;
- ResultsType[2] = N->getValueType(1);
- ResultsType[3] = N->getValueType(2);
-
- NewOps[0] = Chain;
- NewOps[1] = Reg;
- NewOps[2] = Glue;
+ SmallVector<EVT, 4> ResultsType = {MVT::i64, MVT::i64, N->getValueType(1),
+ N->getValueType(2)};
+ SmallVector<SDValue, 3> NewOps = {Chain, Reg, Glue};
SDValue NewValue = DAG.getNode(ISD::CopyFromReg, DL, ResultsType, NewOps);
SDValue Pair = DAG.getNode(ISD::BUILD_PAIR, DL, MVT::i128,
diff --git a/llvm/test/CodeGen/NVPTX/inline-asm-b128-test1.ll b/llvm/test/CodeGen/NVPTX/inline-asm-b128-test1.ll
index 8a256d50d6050..a04ed40dbf91a 100644
--- a/llvm/test/CodeGen/NVPTX/inline-asm-b128-test1.ll
+++ b/llvm/test/CodeGen/NVPTX/inline-asm-b128-test1.ll
@@ -1,3 +1,4 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --extra_scrub --version 5
; RUN: llc < %s -march=nvptx -mcpu=sm_70 -mattr=+ptx83 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_70 -mattr=+ptx83 | %ptxas-verify -arch=sm_70 %}
@@ -5,25 +6,49 @@ target triple = "nvptx64-nvidia-cuda"
@value = internal addrspace(1) global i128 0, align 16
-; Function Attrs: alwaysinline convergent mustprogress willreturn
define void @test_b128_input_from_const() {
- ; CHECK-LABEL: test_b128_input_from_const
- ; CHECK: mov.u64 [[REG_HI:%rd[0-9]+]], 0;
- ; CHECK: mov.u64 [[REG_LO:%rd[0-9]+]], 42;
- ; CHECK: mov.b128 [[REG_128:%rq[0-9]+]], {[[REG_LO]], [[REG_HI]]};
- ; CHECK: { st.b128 [{{%rd[0-9]+}}], [[REG_128]]; }
+; CHECK-LABEL: test_b128_input_from_const(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-NEXT: .reg .b128 %rq<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: mov.u64 %rd2, 0;
+; CHECK-NEXT: mov.u64 %rd3, 42;
+; CHECK-NEXT: mov.b128 %rq1, {%rd3, %rd2};
+; CHECK-NEXT: mov.u32 %r1, value;
+; CHECK-NEXT: cvta.global.u32 %r2, %r1;
+; CHECK-NEXT: cvt.u64.u32 %rd1, %r2;
+; CHECK-NEXT: // begin inline asm
+; CHECK-NEXT: { st.b128 [%rd1], %rq1; }
+; CHECK-NEXT: // end inline asm
+; CHECK-NEXT: ret;
tail call void asm sideeffect "{ st.b128 [$0], $1; }", "l,q"(ptr nonnull addrspacecast (ptr addrspace(1) @value to ptr), i128 42)
ret void
}
-; Function Attrs: alwaysinline convergent mustprogress willreturn
define void @test_b128_input_from_load(ptr nocapture readonly %data) {
- ; CHECK-LABEL: test_b128_input_from_load
- ; CHECK: ld.global.u64 [[REG_HI:%rd[0-9]+]], [[[REG_Addr:%r[0-9]+]]+8];
- ; CHECK: ld.global.u64 [[REG_LO:%rd[0-9]+]], [[[REG_Addr]]];
- ; CHECK: mov.b128 [[REG_128:%rq[0-9]+]], {[[REG_LO]], [[REG_HI]]};
- ; CHECK: { st.b128 [{{%rd[0-9]+}}], [[REG_128]]; }
+; CHECK-LABEL: test_b128_input_from_load(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-NEXT: .reg .b128 %rq<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u32 %r1, [test_b128_input_from_load_param_0];
+; CHECK-NEXT: cvta.to.global.u32 %r2, %r1;
+; CHECK-NEXT: ld.global.u64 %rd2, [%r2+8];
+; CHECK-NEXT: ld.global.u64 %rd3, [%r2];
+; CHECK-NEXT: mov.b128 %rq1, {%rd3, %rd2};
+; CHECK-NEXT: mov.u32 %r3, value;
+; CHECK-NEXT: cvta.global.u32 %r4, %r3;
+; CHECK-NEXT: cvt.u64.u32 %rd1, %r4;
+; CHECK-NEXT: // begin inline asm
+; CHECK-NEXT: { st.b128 [%rd1], %rq1; }
+; CHECK-NEXT: // end inline asm
+; CHECK-NEXT: ret;
%1 = addrspacecast ptr %data to ptr addrspace(1)
%2 = load <2 x i64>, ptr addrspace(1) %1, align 16
@@ -32,13 +57,30 @@ define void @test_b128_input_from_load(ptr nocapture readonly %data) {
ret void
}
-; Function Attrs: alwaysinline convergent mustprogress willreturn
define void @test_b128_input_from_select(ptr nocapture readonly %flag) {
- ; CHECK-LABEL: test_b128_input_from_select
- ; CHECK: selp.b64 [[REG_LO:%rd[0-9]+]], 24, 42, {{%p[0-9]+}};
- ; CHECK: mov.u64 [[REG_HI:%rd[0-9]+]], 0;
- ; CHECK: mov.b128 [[REG_128:%rq[0-9]+]], {[[REG_LO]], [[REG_HI]]};
- ; CHECK: { st.b128 [{{%rd[0-9]+}}], [[REG_128]]; }
+; CHECK-LABEL: test_b128_input_from_select(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<2>;
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-NEXT: .reg .b128 %rq<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u32 %r1, [test_b128_input_from_select_param_0];
+; CHECK-NEXT: cvta.to.global.u32 %r2, %r1;
+; CHECK-NEXT: ld.global.u8 %rs1, [%r2];
+; CHECK-NEXT: setp.eq.s16 %p1, %rs1, 0;
+; CHECK-NEXT: selp.b64 %rd2, 24, 42, %p1;
+; CHECK-NEXT: mov.u64 %rd3, 0;
+; CHECK-NEXT: mov.b128 %rq1, {%rd2, %rd3};
+; CHECK-NEXT: mov.u32 %r3, value;
+; CHECK-NEXT: cvta.global.u32 %r4, %r3;
+; CHECK-NEXT: cvt.u64.u32 %rd1, %r4;
+; CHECK-NEXT: // begin inline asm
+; CHECK-NEXT: { st.b128 [%rd1], %rq1; }
+; CHECK-NEXT: // end inline asm
+; CHECK-NEXT: ret;
%1 = addrspacecast ptr %flag to ptr addrspace(1)
%tmp1 = load i8, ptr addrspace(1) %1, align 1
@@ -48,12 +90,23 @@ define void @test_b128_input_from_select(ptr nocapture readonly %flag) {
ret void
}
-; Function Attrs: alwaysinline mustprogress willreturn memory(write, argmem: none, inaccessiblemem: none)
define void @test_store_b128_output() {
- ; CHECK-LABEL: test_store_b128_output
- ; CHECK: { mov.b128 [[REG_128:%rq[0-9]+]], 41; }
- ; CHECK: mov.b128 {%rd{{[0-9]+}}, %rd{{[0-9]+}}}, [[REG_128]];
-
+; CHECK-LABEL: test_store_b128_output(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<5>;
+; CHECK-NEXT: .reg .b128 %rq<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: // begin inline asm
+; CHECK-NEXT: { mov.b128 %rq1, 41; }
+; CHECK-NEXT: // end inline asm
+; CHECK-NEXT: mov.b128 {%rd1, %rd2}, %rq1;
+; CHECK-NEXT: add.cc.s64 %rd3, %rd1, 1;
+; CHECK-NEXT: addc.cc.s64 %rd4, %rd2, 0;
+; CHECK-NEXT: st.global.u64 [value+8], %rd4;
+; CHECK-NEXT: st.global.u64 [value], %rd3;
+; CHECK-NEXT: ret;
+
%1 = tail call i128 asm "{ mov.b128 $0, 41; }", "=q"()
%add = add nsw i128 %1, 1
%2 = bitcast i128 %add to <2 x i64>
@@ -61,14 +114,28 @@ define void @test_store_b128_output() {
ret void
}
-; Function Attrs: alwaysinline mustprogress willreturn memory(write, argmem: read, inaccessiblemem: none)
define void @test_use_of_b128_output(ptr nocapture readonly %data) {
- ; CHECK-LABEL: test_use_of_b128_output
- ; CHECK: ld.global.u64 [[REG_HI:%rd[0-9]+]], [[[REG_Addr:%r[0-9]+]]+8];
- ; CHECK: ld.global.u64 [[REG_LO:%rd[0-9]+]], [[[REG_Addr]]];
- ; CHECK: mov.b128 [[REG_128_IN:%rq[0-9]+]], {[[REG_LO]], [[REG_HI]]};
- ; CHECK: { mov.b128 [[REG_128_OUT:%rq[0-9]+]], [[REG_128_IN]]; }
- ; CHECK: mov.b128 {%rd{{[0-9]+}}, %rd{{[0-9]+}}}, [[REG_128_OUT]];
+; CHECK-LABEL: test_use_of_b128_output(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<7>;
+; CHECK-NEXT: .reg .b128 %rq<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u32 %r1, [test_use_of_b128_output_param_0];
+; CHECK-NEXT: cvta.to.global.u32 %r2, %r1;
+; CHECK-NEXT: ld.global.u64 %rd1, [%r2+8];
+; CHECK-NEXT: ld.global.u64 %rd2, [%r2];
+; CHECK-NEXT: mov.b128 %rq2, {%rd2, %rd1};
+; CHECK-NEXT: // begin inline asm
+; CHECK-NEXT: { mov.b128 %rq1, %rq2; }
+; CHECK-NEXT: // end inline asm
+; CHECK-NEXT: mov.b128 {%rd3, %rd4}, %rq1;
+; CHECK-NEXT: add.cc.s64 %rd5, %rd3, 1;
+; CHECK-NEXT: addc.cc.s64 %rd6, %rd4, 0;
+; CHECK-NEXT: st.global.u64 [value], %rd5;
+; CHECK-NEXT: st.global.u64 [value+8], %rd6;
+; CHECK-NEXT: ret;
%1 = addrspacecast ptr %data to ptr addrspace(1)
%2 = load <2 x i64>, ptr addrspace(1) %1, align 16
diff --git a/llvm/test/CodeGen/NVPTX/inline-asm-b128-test2.ll b/llvm/test/CodeGen/NVPTX/inline-asm-b128-test2.ll
index 09b648d036a4c..bb45ff6ba2e27 100644
--- a/llvm/test/CodeGen/NVPTX/inline-asm-b128-test2.ll
+++ b/llvm/test/CodeGen/NVPTX/inline-asm-b128-test2.ll
@@ -1,3 +1,4 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --extra_scrub --version 5
; RUN: llc < %s -march=nvptx -mcpu=sm_70 -mattr=+ptx83 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_70 -mattr=+ptx83 | %ptxas-verify -arch=sm_70 %}
@@ -13,33 +14,110 @@ target triple = "nvptx64-nvidia-cuda"
@v_i128_min = internal addrspace(1) global i128 0, align 16
@v64 = internal addrspace(1) global ptr null, align 8
-; Function Attrs: alwaysinline convergent mustprogress willreturn
define void @test_corner_values() {
- ; CHECK-LABEL: test_corner_values
- ; CHECK: mov.u64 [[U64_MAX:%rd[0-9]+]], -1;
- ; CHECK: mov.b128 [[U128_MAX:%rq[0-9]+]], {[[U64_MAX]], [[U64_MAX]]};
- ; CHECK: mov.u64 [[I64_MAX:%rd[0-9]+]], 9223372036854775807;
- ; CHECK: mov.b128 [[I128_MAX:%rq[0-9]+]], {[[U64_MAX]], [[I64_MAX]]}
- ; CHECK: mov.u64 [[I64_MIN:%rd[0-9]+]], -9223372036854775808;
- ; CHECK: mov.u64 [[U64_ZERO:%rd[0-9]+]], 0;
- ; CHECK: mov.b128 [[I128_MIN:%rq[0-9]+]], {[[U64_ZERO]], [[I64_MIN]]}
- ; CHECK: mov.b128 [[U128_ZERO:%rq[0-9]+]], {[[U64_ZERO]], [[U64_ZERO]]}
+; CHECK-LABEL: test_corner_values(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<20>;
+; CHECK-NEXT: .reg .b64 %rd<17>;
+; CHECK-NEXT: .reg .b128 %rq<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.global.u32 %r1, [v64];
+; CHECK-NEXT: add.s32 %r2, %r1, 8;
+; CHECK-NEXT: mov.u64 %rd13, -1;
+; CHECK-NEXT: mov.b128 %rq1, {%rd13, %rd13};
+; CHECK-NEXT: cvt.u64.u32 %rd1, %r1;
+; CHECK-NEXT: cvt.u64.u32 %rd2, %r2;
+; CHECK-NEXT: mov.u32 %r3, v_u128_max;
+; CHECK-NEXT: cvta.global.u32 %r4, %r3;
+; CHECK-NEXT: cvt.u64.u32 %rd3, %r4;
+; CHECK-NEXT: // begin inline asm
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b64 hi;
+; CHECK-NEXT: .reg .b64 lo;
+; CHECK-NEXT: mov.b128 {lo, hi}, %rq1;
+; CHECK-NEXT: st.b64 [%rd1], lo;
+; CHECK-NEXT: st.b64 [%rd2], hi;
+; CHECK-NEXT: st.b128 [%rd3], %rq1;
+; CHECK-NEXT: }
+; CHECK-NEXT: // end inline asm
+; CHECK-NEXT: ld.global.u32 %r5, [v64];
+; CHECK-NEXT: add.s32 %r6, %r5, 16;
+; CHECK-NEXT: add.s32 %r7, %r5, 24;
+; CHECK-NEXT: mov.u64 %rd14, 9223372036854775807;
+; CHECK-NEXT: mov.b128 %rq2, {%rd13, %rd14};
+; CHECK-NEXT: mov.u32 %r8, v_i128_max;
+; CHECK-NEXT: cvta.global.u32 %r9, %r8;
+; CHECK-NEXT: cvt.u64.u32 %rd6, %r9;
+; CHECK-NEXT: cvt.u64.u32 %rd4, %r6;
+; CHECK-NEXT: cvt.u64.u32 %rd5, %r7;
+; CHECK-NEXT: // begin inline asm
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b64 hi;
+; CHECK-NEXT: .reg .b64 lo;
+; CHECK-NEXT: mov.b128 {lo, hi}, %rq2;
+; CHECK-NEXT: st.b64 [%rd4], lo;
+; CHECK-NEXT: st.b64 [%rd5], hi;
+; CHECK-NEXT: st.b128 [%rd6], %rq2;
+; CHECK-NEXT: }
+; CHECK-NEXT: // end inline asm
+; CHECK-NEXT: ld.global.u32 %r10, [v64];
+; CHECK-NEXT: add.s32 %r11, %r10, 32;
+; CHECK-NEXT: add.s32 %r12, %r10, 40;
+; CHECK-NEXT: mov.u64 %rd15, -9223372036854775808;
+; CHECK-NEXT: mov.u64 %rd16, 0;
+; CHECK-NEXT: mov.b128 %rq3, {%rd16, %rd15};
+; CHECK-NEXT: mov.u32 %r13, v_i128_min;
+; CHECK-NEXT: cvta.global.u32 %r14, %r13;
+; CHECK-NEXT: cvt.u64.u32 %rd9, %r14;
+; CHECK-NEXT: cvt.u64.u32 %rd7, %r11;
+; CHECK-NEXT: cvt.u64.u32 %rd8, %r12;
+; CHECK-NEXT: // begin inline asm
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b64 hi;
+; CHECK-NEXT: .reg .b64 lo;
+; CHECK-NEXT: mov.b128 {lo, hi}, %rq3;
+; CHECK-NEXT: st.b64 [%rd7], lo;
+; CHECK-NEXT: st.b64 [%rd8], hi;
+; CHECK-NEXT: st.b128 [%rd9], %rq3;
+; CHECK-NEXT: }
+; CHECK-NEXT: // end inline asm
+; CHECK-NEXT: ld.global.u32 %r15, [v64];
+; CHECK-NEXT: add.s32 %r16, %r15, 48;
+; CHECK-NEXT: add.s32 %r17, %r15, 56;
+; CHECK-NEXT: mov.b128 %rq4, {%rd16, %rd16};
+; CHECK-NEXT: mov.u32 %r18, v_u128_zero;
+; CHECK-NEXT: cvta.global.u32 %r19, %r18;
+; CHECK-NEXT: cvt.u64.u32 %rd12, %r19;
+; CHECK-NEXT: cvt.u64.u32 %rd10, %r16;
+; CHECK-NEXT: cvt.u64.u32 %rd11, %r17;
+; CHECK-NEXT: // begin inline asm
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b64 hi;
+; CHECK-NEXT: .reg .b64 lo;
+; CHECK-NEXT: mov.b128 {lo, hi}, %rq4;
+; CHECK-NEXT: st.b64 [%rd10], lo;
+; CHECK-NEXT: st.b64 [%rd11], hi;
+; CHECK-NEXT: st.b128 [%rd12], %rq4;
+; CHECK-NEXT: }
+; CHECK-NEXT: // end inline asm
+; CHECK-NEXT: ret;
%tmp = load ptr, ptr addrspace(1) @v64, align 8
%add.ptr2 = getelementptr inbounds i64, ptr %tmp, i64 1
- tail call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}\0A\09", "q,l,l,l"(i128 -1, ptr %tmp, ptr nonnull %add.ptr2, ptr nonnull addrspacecast (ptr addrspace(1) @v_u128_max to ptr))
+ tail call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}", "q,l,l,l"(i128 -1, ptr %tmp, ptr nonnull %add.ptr2, ptr nonnull addrspacecast (ptr addrspace(1) @v_u128_max to ptr))
%tmp3 = load ptr, ptr addrspace(1) @v64, align 8
%add.ptr4 = getelementptr inbounds i64, ptr %tmp3, i64 2
%add.ptr6 = getelementptr inbounds i64, ptr %tmp3, i64 3
- tail call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}\0A\09", "q,l,l,l"(i128 170141183460469231731687303715884105727, ptr nonnull %add.ptr4, ptr nonnull %add.ptr6, ptr nonnull addrspacecast (ptr addrspace(1) @v_i128_max to ptr))
+ tail call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}", "q,l,l,l"(i128 170141183460469231731687303715884105727, ptr nonnull %add.ptr4, ptr nonnull %add.ptr6, ptr nonnull addrspacecast (ptr addrspace(1) @v_i128_max to ptr))
%tmp7 = load ptr, ptr addrspace(1) @v64, align 8
%add.ptr8 = getelementptr inbounds i64, ptr %tmp7, i64 4
%add.ptr10 = getelementptr inbounds i64, ptr %tmp7, i64 5
- tail call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}\0A\09", "q,l,l,l"(i128 -170141183460469231731687303715884105728, ptr nonnull %add.ptr8, ptr nonnull %add.ptr10, ptr nonnull addrspacecast (ptr addrspace(1) @v_i128_min to ptr))
+ tail call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}", "q,l,l,l"(i128 -170141183460469231731687303715884105728, ptr nonnull %add.ptr8, ptr nonnull %add.ptr10, ptr nonnull addrspacecast (ptr addrspace(1) @v_i128_min to ptr))
%tmp11 = load ptr, ptr addrspace(1) @v64, align 8
%add.ptr12 = getelementptr inbounds i64, ptr %tmp11, i64 6
%add.ptr14 = getelementptr inbounds i64, ptr %tmp11, i64 7
- tail call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}\0A\09", "q,l,l,l"(i128 0, ptr nonnull %add.ptr12, ptr nonnull %add.ptr14, ptr nonnull addrspacecast (ptr addrspace(1) @v_u128_zero to ptr))
+ tail call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}", "q,l,l,l"(i128 0, ptr nonnull %add.ptr12, ptr nonnull %add.ptr14, ptr nonnull addrspacecast (ptr addrspace(1) @v_u128_zero to ptr))
ret void
}
diff --git a/llvm/test/CodeGen/NVPTX/inline-asm-b128-test3.ll b/llvm/test/CodeGen/NVPTX/inline-asm-b128-test3.ll
index e187aa4370858..081956447345c 100644
--- a/llvm/test/CodeGen/NVPTX/inline-asm-b128-test3.ll
+++ b/llvm/test/CodeGen/NVPTX/inline-asm-b128-test3.ll
@@ -1,3 +1,4 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --extra_scrub --version 5
; RUN: llc < %s -march=nvptx -mcpu=sm_70 -mattr=+ptx83 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_70 -mattr=+ptx83 | %ptxas-verify -arch=sm_70 %}
@@ -7,13 +8,40 @@ target triple = "nvptx64-nvidia-cuda"
@x = internal addrspace(1) global i128 0, align 16
define void @test_b128_in_loop() {
- ; CHECK-LABEL: test_b128_in_loop
- ; CHECK: ld.global.u64 [[REG_HI:%rd[0-9]+]], [x+8];
- ; CHECK: ld.global.u64 [[REG_LO:%rd[0-9]+]], [x];
- ; CHECK: mov.b128 [[REG_128:%rq[0-9]+]], {[[REG_LO]], [[REG_HI]]};
- ; CHECK: mov.b128 {lo, hi}, [[REG_128]];
- ; CHECK: add.cc.u64 lo, lo, {{%rd[0-9]+}};
- ; CHECK: mov.b128 [[REG_128]], {lo, hi};
+; CHECK-LABEL: test_b128_in_loop(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<3>;
+; CHECK-NEXT: .reg .b64 %rd<15>;
+; CHECK-NEXT: .reg .b128 %rq<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.global.s32 %rd1, [size];
+; CHECK-NEXT: setp.eq.s64 %p1, %rd1, 0;
+; CHECK-NEXT: @%p1 bra $L__BB0_3;
+; CHECK-NEXT: // %bb.1: // %.lr.ph.preheader
+; CHECK-NEXT: ld.global.u64 %rd13, [x+8];
+; CHECK-NEXT: ld.global.u64 %rd12, [x];
+; CHECK-NEXT: mov.u64 %rd14, 0;
+; CHECK-NEXT: $L__BB0_2: // %.lr.ph
+; CHECK-NEXT: // =>This Inner Loop Header: Depth=1
+; CHECK-NEXT: mov.b128 %rq1, {%rd12, %rd13};
+; CHECK-NEXT: // begin inline asm
+; CHECK-NEXT: {
+; CHECK-NEXT: .reg .b64 lo;
+; CHECK-NEXT: .reg .b64 hi;
+; CHECK-NEXT: mov.b128 {lo, hi}, %rq1;
+; CHECK-NEXT: add.cc.u64 lo, lo, %rd14;
+; CHECK-NEXT: mov.b128 %rq1, {lo, hi};
+; CHECK-NEXT: }
+; CHECK-NEXT: // end inline asm
+; CHECK-NEXT: mov.b128 {%rd12, %rd13}, %rq1;
+; CHECK-NEXT: st.global.u64 [x+8], %rd13;
+; CHECK-NEXT: st.global.u64 [x], %rd12;
+; CHECK-NEXT: add.s64 %rd14, %rd14, 1;
+; CHECK-NEXT: setp.ne.s64 %p2, %rd1, %rd14;
+; CHECK-NEXT: @%p2 bra $L__BB0_2;
+; CHECK-NEXT: $L__BB0_3: // %._crit_edge
+; CHECK-NEXT: ret;
%tmp11 = load i32, ptr addrspace(1) @size, align 4
%cmp3.not = icmp eq i32 %tmp11, 0
@@ -27,7 +55,7 @@ define void @test_b128_in_loop() {
.lr.ph: ; preds = %.lr.ph, %.lr.ph.preheader
%1 = phi i128 [ %2, %.lr.ph ], [ %x.promoted5, %.lr.ph.preheader ]
%i.04 = phi i64 [ %inc, %.lr.ph ], [ 0, %.lr.ph.preheader ]
- %2 = tail call i128 asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09add.cc.u64 lo, lo, $1;\0A\09mov.b128 $0, {lo, hi};\0A\09} \0A\09", "=q,l,0"(i64 %i.04, i128 %1)
+ %2 = tail call i128 asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09add.cc.u64 lo, lo, $1;\0A\09mov.b128 $0, {lo, hi};\0A\09}", "=q,l,0"(i64 %i.04, i128 %1)
%3 = bitcast i128 %2 to <2 x i64>
store <2 x i64> %3, ptr addrspace(1) @x, align 16
%inc = add nuw i64 %i.04, 1
>From 6bb8dd74a609120c0b082161a575d4d5ba39ee32 Mon Sep 17 00:00:00 2001
From: chengjunp <chengjunp at nvidia.com>
Date: Wed, 26 Jun 2024 22:25:04 +0000
Subject: [PATCH 5/6] Rename variables in tests & Use dag helper to build nodes
---
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 4 +-
.../CodeGen/NVPTX/inline-asm-b128-test1.ll | 13 ++----
.../CodeGen/NVPTX/inline-asm-b128-test2.ll | 36 +++++++--------
.../CodeGen/NVPTX/inline-asm-b128-test3.ll | 45 +++++++++----------
4 files changed, 41 insertions(+), 57 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index fd4293e56cfb6..c02d874a9a6b3 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -3111,14 +3111,14 @@ SDValue NVPTXTargetLowering::LowerCopyToReg_128(SDValue Op,
SDNode *Node = Op.getNode();
SDLoc DL(Node);
- SDValue Cast = DAG.getNode(ISD::BITCAST, DL, MVT::v2i64, Op->getOperand(2));
+ SDValue Cast = DAG.getBitcast(MVT::v2i64, Op->getOperand(2));
SDValue Lo = DAG.getNode(ISD::EXTRACT_VECTOR_ELT, DL, MVT::i64, Cast,
DAG.getIntPtrConstant(0, DL));
SDValue Hi = DAG.getNode(ISD::EXTRACT_VECTOR_ELT, DL, MVT::i64, Cast,
DAG.getIntPtrConstant(1, DL));
SmallVector<SDValue, 5> NewOps(Op->getNumOperands() + 1);
- SmallVector<EVT, 3> ResultsType(Node->value_begin(), Node->value_end());
+ SmallVector<EVT, 3> ResultsType(Node->values());
NewOps[0] = Op->getOperand(0); // Chain
NewOps[1] = Op->getOperand(1); // Dst Reg
diff --git a/llvm/test/CodeGen/NVPTX/inline-asm-b128-test1.ll b/llvm/test/CodeGen/NVPTX/inline-asm-b128-test1.ll
index a04ed40dbf91a..3232f40a40a70 100644
--- a/llvm/test/CodeGen/NVPTX/inline-asm-b128-test1.ll
+++ b/llvm/test/CodeGen/NVPTX/inline-asm-b128-test1.ll
@@ -83,10 +83,10 @@ define void @test_b128_input_from_select(ptr nocapture readonly %flag) {
; CHECK-NEXT: ret;
%1 = addrspacecast ptr %flag to ptr addrspace(1)
- %tmp1 = load i8, ptr addrspace(1) %1, align 1
- %tobool.not = icmp eq i8 %tmp1, 0
- %. = select i1 %tobool.not, i128 24, i128 42
- tail call void asm sideeffect "{ st.b128 [$0], $1; }", "l,q"(ptr nonnull addrspacecast (ptr addrspace(1) @value to ptr), i128 %.)
+ %2 = load i8, ptr addrspace(1) %1, align 1
+ %3 = icmp eq i8 %2, 0
+ %4 = select i1 %3, i128 24, i128 42
+ tail call void asm sideeffect "{ st.b128 [$0], $1; }", "l,q"(ptr nonnull addrspacecast (ptr addrspace(1) @value to ptr), i128 %4)
ret void
}
@@ -146,8 +146,3 @@ define void @test_use_of_b128_output(ptr nocapture readonly %data) {
store <2 x i64> %5, ptr addrspace(1) @value, align 16
ret void
}
-
-!nvvmir.version = !{!0, !1, !0, !1, !1, !0, !0, !0, !1}
-
-!0 = !{i32 2, i32 0, i32 3, i32 1}
-!1 = !{i32 2, i32 0}
diff --git a/llvm/test/CodeGen/NVPTX/inline-asm-b128-test2.ll b/llvm/test/CodeGen/NVPTX/inline-asm-b128-test2.ll
index bb45ff6ba2e27..3d1d7fbbe27e8 100644
--- a/llvm/test/CodeGen/NVPTX/inline-asm-b128-test2.ll
+++ b/llvm/test/CodeGen/NVPTX/inline-asm-b128-test2.ll
@@ -103,26 +103,20 @@ define void @test_corner_values() {
; CHECK-NEXT: // end inline asm
; CHECK-NEXT: ret;
- %tmp = load ptr, ptr addrspace(1) @v64, align 8
- %add.ptr2 = getelementptr inbounds i64, ptr %tmp, i64 1
- tail call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}", "q,l,l,l"(i128 -1, ptr %tmp, ptr nonnull %add.ptr2, ptr nonnull addrspacecast (ptr addrspace(1) @v_u128_max to ptr))
- %tmp3 = load ptr, ptr addrspace(1) @v64, align 8
- %add.ptr4 = getelementptr inbounds i64, ptr %tmp3, i64 2
- %add.ptr6 = getelementptr inbounds i64, ptr %tmp3, i64 3
- tail call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}", "q,l,l,l"(i128 170141183460469231731687303715884105727, ptr nonnull %add.ptr4, ptr nonnull %add.ptr6, ptr nonnull addrspacecast (ptr addrspace(1) @v_i128_max to ptr))
- %tmp7 = load ptr, ptr addrspace(1) @v64, align 8
- %add.ptr8 = getelementptr inbounds i64, ptr %tmp7, i64 4
- %add.ptr10 = getelementptr inbounds i64, ptr %tmp7, i64 5
- tail call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}", "q,l,l,l"(i128 -170141183460469231731687303715884105728, ptr nonnull %add.ptr8, ptr nonnull %add.ptr10, ptr nonnull addrspacecast (ptr addrspace(1) @v_i128_min to ptr))
- %tmp11 = load ptr, ptr addrspace(1) @v64, align 8
- %add.ptr12 = getelementptr inbounds i64, ptr %tmp11, i64 6
- %add.ptr14 = getelementptr inbounds i64, ptr %tmp11, i64 7
- tail call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}", "q,l,l,l"(i128 0, ptr nonnull %add.ptr12, ptr nonnull %add.ptr14, ptr nonnull addrspacecast (ptr addrspace(1) @v_u128_zero to ptr))
+ %1 = load ptr, ptr addrspace(1) @v64, align 8
+ %2 = getelementptr inbounds i64, ptr %1, i64 1
+ tail call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}", "q,l,l,l"(i128 -1, ptr %1, ptr nonnull %2, ptr nonnull addrspacecast (ptr addrspace(1) @v_u128_max to ptr))
+ %3 = load ptr, ptr addrspace(1) @v64, align 8
+ %4 = getelementptr inbounds i64, ptr %3, i64 2
+ %5 = getelementptr inbounds i64, ptr %3, i64 3
+ tail call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}", "q,l,l,l"(i128 170141183460469231731687303715884105727, ptr nonnull %4, ptr nonnull %5, ptr nonnull addrspacecast (ptr addrspace(1) @v_i128_max to ptr))
+ %6 = load ptr, ptr addrspace(1) @v64, align 8
+ %7 = getelementptr inbounds i64, ptr %6, i64 4
+ %8 = getelementptr inbounds i64, ptr %6, i64 5
+ tail call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}", "q,l,l,l"(i128 -170141183460469231731687303715884105728, ptr nonnull %7, ptr nonnull %8, ptr nonnull addrspacecast (ptr addrspace(1) @v_i128_min to ptr))
+ %9 = load ptr, ptr addrspace(1) @v64, align 8
+ %10 = getelementptr inbounds i64, ptr %9, i64 6
+ %11 = getelementptr inbounds i64, ptr %9, i64 7
+ tail call void asm sideeffect "{\0A\09.reg .b64 hi;\0A\09.reg .b64 lo;\0A\09mov.b128 {lo, hi}, $0;\0A\09st.b64 [$1], lo;\0A\09st.b64 [$2], hi;\0A\09st.b128 [$3], $0;\0A\09}", "q,l,l,l"(i128 0, ptr nonnull %10, ptr nonnull %11, ptr nonnull addrspacecast (ptr addrspace(1) @v_u128_zero to ptr))
ret void
}
-
-
-!nvvmir.version = !{!2, !3, !2, !3, !3, !2, !2, !2, !3}
-
-!2 = !{i32 2, i32 0, i32 3, i32 1}
-!3 = !{i32 2, i32 0}
diff --git a/llvm/test/CodeGen/NVPTX/inline-asm-b128-test3.ll b/llvm/test/CodeGen/NVPTX/inline-asm-b128-test3.ll
index 081956447345c..ae453977123e0 100644
--- a/llvm/test/CodeGen/NVPTX/inline-asm-b128-test3.ll
+++ b/llvm/test/CodeGen/NVPTX/inline-asm-b128-test3.ll
@@ -18,11 +18,11 @@ define void @test_b128_in_loop() {
; CHECK-NEXT: ld.global.s32 %rd1, [size];
; CHECK-NEXT: setp.eq.s64 %p1, %rd1, 0;
; CHECK-NEXT: @%p1 bra $L__BB0_3;
-; CHECK-NEXT: // %bb.1: // %.lr.ph.preheader
+; CHECK-NEXT: // %bb.1: // %BB1
; CHECK-NEXT: ld.global.u64 %rd13, [x+8];
; CHECK-NEXT: ld.global.u64 %rd12, [x];
; CHECK-NEXT: mov.u64 %rd14, 0;
-; CHECK-NEXT: $L__BB0_2: // %.lr.ph
+; CHECK-NEXT: $L__BB0_2: // %BB2
; CHECK-NEXT: // =>This Inner Loop Header: Depth=1
; CHECK-NEXT: mov.b128 %rq1, {%rd12, %rd13};
; CHECK-NEXT: // begin inline asm
@@ -40,33 +40,28 @@ define void @test_b128_in_loop() {
; CHECK-NEXT: add.s64 %rd14, %rd14, 1;
; CHECK-NEXT: setp.ne.s64 %p2, %rd1, %rd14;
; CHECK-NEXT: @%p2 bra $L__BB0_2;
-; CHECK-NEXT: $L__BB0_3: // %._crit_edge
+; CHECK-NEXT: $L__BB0_3: // %BB3
; CHECK-NEXT: ret;
- %tmp11 = load i32, ptr addrspace(1) @size, align 4
- %cmp3.not = icmp eq i32 %tmp11, 0
- br i1 %cmp3.not, label %._crit_edge, label %.lr.ph.preheader
+ %1 = load i32, ptr addrspace(1) @size, align 4
+ %2 = icmp eq i32 %1, 0
+ br i1 %2, label %BB3, label %BB1
-.lr.ph.preheader: ; preds = %0
- %x.promoted5 = load i128, ptr addrspace(1) @x, align 16
- %umax = sext i32 %tmp11 to i64
- br label %.lr.ph
+BB1: ; preds = %0
+ %3 = load i128, ptr addrspace(1) @x, align 16
+ %4 = sext i32 %1 to i64
+ br label %BB2
-.lr.ph: ; preds = %.lr.ph, %.lr.ph.preheader
- %1 = phi i128 [ %2, %.lr.ph ], [ %x.promoted5, %.lr.ph.preheader ]
- %i.04 = phi i64 [ %inc, %.lr.ph ], [ 0, %.lr.ph.preheader ]
- %2 = tail call i128 asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09add.cc.u64 lo, lo, $1;\0A\09mov.b128 $0, {lo, hi};\0A\09}", "=q,l,0"(i64 %i.04, i128 %1)
- %3 = bitcast i128 %2 to <2 x i64>
- store <2 x i64> %3, ptr addrspace(1) @x, align 16
- %inc = add nuw i64 %i.04, 1
- %exitcond.not = icmp eq i64 %inc, %umax
- br i1 %exitcond.not, label %._crit_edge, label %.lr.ph
+BB2: ; preds = %BB2, %BB1
+ %5 = phi i128 [ %7, %BB2 ], [ %3, %BB1 ]
+ %6 = phi i64 [ %9, %BB2 ], [ 0, %BB1 ]
+ %7 = tail call i128 asm "{\0A\09.reg .b64 lo;\0A\09.reg .b64 hi;\0A\09mov.b128 {lo, hi}, $0;\0A\09add.cc.u64 lo, lo, $1;\0A\09mov.b128 $0, {lo, hi};\0A\09}", "=q,l,0"(i64 %6, i128 %5)
+ %8 = bitcast i128 %7 to <2 x i64>
+ store <2 x i64> %8, ptr addrspace(1) @x, align 16
+ %9 = add nuw i64 %6, 1
+ %10 = icmp eq i64 %9, %4
+ br i1 %10, label %BB3, label %BB2
-._crit_edge: ; preds = %.lr.ph, %0
+BB3: ; preds = %BB2, %0
ret void
}
-
-!nvvmir.version = !{!0, !1, !0, !1, !1, !0, !0, !0, !1}
-
-!0 = !{i32 2, i32 0, i32 3, i32 1}
-!1 = !{i32 2, i32 0}
>From e7a3af718bfdbd703a4f616977303649d4b88209 Mon Sep 17 00:00:00 2001
From: chengjunp <chengjunp at nvidia.com>
Date: Fri, 28 Jun 2024 22:05:15 +0000
Subject: [PATCH 6/6] Update LangRef for the new constraint code for 128-bit
values in NVPTX.
---
llvm/docs/LangRef.rst | 1 +
1 file changed, 1 insertion(+)
diff --git a/llvm/docs/LangRef.rst b/llvm/docs/LangRef.rst
index edb362c617565..d5d5a4622da4c 100644
--- a/llvm/docs/LangRef.rst
+++ b/llvm/docs/LangRef.rst
@@ -5381,6 +5381,7 @@ NVPTX:
- ``c`` or ``h``: A 16-bit integer register.
- ``r``: A 32-bit integer register.
- ``l`` or ``N``: A 64-bit integer register.
+- ``q``: A 128-bit integer register.
- ``f``: A 32-bit float register.
- ``d``: A 64-bit float register.
More information about the cfe-commits
mailing list