[llvm] d3dae84 - [NVPTX] Switch to imm offset variants for LDG and LDU (#128270)
via llvm-commits
llvm-commits at lists.llvm.org
Mon Feb 24 23:58:48 PST 2025
Author: Alex MacLean
Date: 2025-02-24T23:58:45-08:00
New Revision: d3dae841c05c9447b665a8334aa3cfeac904d749
URL: https://github.com/llvm/llvm-project/commit/d3dae841c05c9447b665a8334aa3cfeac904d749
DIFF: https://github.com/llvm/llvm-project/commit/d3dae841c05c9447b665a8334aa3cfeac904d749.diff
LOG: [NVPTX] Switch to imm offset variants for LDG and LDU (#128270)
Added:
Modified:
llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
llvm/test/CodeGen/NVPTX/ldu-ldg.ll
llvm/test/CodeGen/NVPTX/variadics-backend.ll
Removed:
################################################################################
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
index 6f0bf510ad893..971a128aadfdb 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
@@ -1106,9 +1106,6 @@ bool NVPTXDAGToDAGISel::tryLoadVector(SDNode *N) {
std::nullopt, NVPTX::LDV_f32_v4_asi, std::nullopt);
break;
}
- if (!Opcode)
- return false;
- Ops.append({Base, Offset, Chain});
} else {
if (PointerSize == 64) {
SelectADDRri64(Op1.getNode(), Op1, Base, Offset);
@@ -1148,10 +1145,10 @@ bool NVPTXDAGToDAGISel::tryLoadVector(SDNode *N) {
break;
}
}
- if (!Opcode)
- return false;
- Ops.append({Base, Offset, Chain});
}
+ if (!Opcode)
+ return false;
+ Ops.append({Base, Offset, Chain});
LD = CurDAG->getMachineNode(*Opcode, DL, N->getVTList(), Ops);
MachineMemOperand *MemRef = cast<MemSDNode>(N)->getMemOperand();
@@ -1202,63 +1199,59 @@ bool NVPTXDAGToDAGISel::tryLDGLDU(SDNode *N) {
std::optional<unsigned> Opcode;
SDLoc DL(N);
SDNode *LD;
- SDValue Base, Offset, Addr;
+ SDValue Base, Offset;
- if (SelectDirectAddr(Op1, Addr)) {
+ if (SelectADDRsi(Op1.getNode(), Op1, Base, Offset)) {
switch (N->getOpcode()) {
default:
return false;
case ISD::LOAD:
Opcode = pickOpcodeForVT(
- EltVT.getSimpleVT().SimpleTy, NVPTX::INT_PTX_LDG_GLOBAL_i8avar,
- NVPTX::INT_PTX_LDG_GLOBAL_i16avar, NVPTX::INT_PTX_LDG_GLOBAL_i32avar,
- NVPTX::INT_PTX_LDG_GLOBAL_i64avar, NVPTX::INT_PTX_LDG_GLOBAL_f32avar,
- NVPTX::INT_PTX_LDG_GLOBAL_f64avar);
+ EltVT.getSimpleVT().SimpleTy, NVPTX::INT_PTX_LDG_GLOBAL_i8asi,
+ NVPTX::INT_PTX_LDG_GLOBAL_i16asi, NVPTX::INT_PTX_LDG_GLOBAL_i32asi,
+ NVPTX::INT_PTX_LDG_GLOBAL_i64asi, NVPTX::INT_PTX_LDG_GLOBAL_f32asi,
+ NVPTX::INT_PTX_LDG_GLOBAL_f64asi);
break;
case ISD::INTRINSIC_W_CHAIN:
Opcode = pickOpcodeForVT(
- EltVT.getSimpleVT().SimpleTy, NVPTX::INT_PTX_LDU_GLOBAL_i8avar,
- NVPTX::INT_PTX_LDU_GLOBAL_i16avar, NVPTX::INT_PTX_LDU_GLOBAL_i32avar,
- NVPTX::INT_PTX_LDU_GLOBAL_i64avar, NVPTX::INT_PTX_LDU_GLOBAL_f32avar,
- NVPTX::INT_PTX_LDU_GLOBAL_f64avar);
+ EltVT.getSimpleVT().SimpleTy, NVPTX::INT_PTX_LDU_GLOBAL_i8asi,
+ NVPTX::INT_PTX_LDU_GLOBAL_i16asi, NVPTX::INT_PTX_LDU_GLOBAL_i32asi,
+ NVPTX::INT_PTX_LDU_GLOBAL_i64asi, NVPTX::INT_PTX_LDU_GLOBAL_f32asi,
+ NVPTX::INT_PTX_LDU_GLOBAL_f64asi);
break;
case NVPTXISD::LoadV2:
Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
- NVPTX::INT_PTX_LDG_G_v2i8_ELE_avar,
- NVPTX::INT_PTX_LDG_G_v2i16_ELE_avar,
- NVPTX::INT_PTX_LDG_G_v2i32_ELE_avar,
- NVPTX::INT_PTX_LDG_G_v2i64_ELE_avar,
- NVPTX::INT_PTX_LDG_G_v2f32_ELE_avar,
- NVPTX::INT_PTX_LDG_G_v2f64_ELE_avar);
+ NVPTX::INT_PTX_LDG_G_v2i8_ELE_asi,
+ NVPTX::INT_PTX_LDG_G_v2i16_ELE_asi,
+ NVPTX::INT_PTX_LDG_G_v2i32_ELE_asi,
+ NVPTX::INT_PTX_LDG_G_v2i64_ELE_asi,
+ NVPTX::INT_PTX_LDG_G_v2f32_ELE_asi,
+ NVPTX::INT_PTX_LDG_G_v2f64_ELE_asi);
break;
case NVPTXISD::LDUV2:
Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
- NVPTX::INT_PTX_LDU_G_v2i8_ELE_avar,
- NVPTX::INT_PTX_LDU_G_v2i16_ELE_avar,
- NVPTX::INT_PTX_LDU_G_v2i32_ELE_avar,
- NVPTX::INT_PTX_LDU_G_v2i64_ELE_avar,
- NVPTX::INT_PTX_LDU_G_v2f32_ELE_avar,
- NVPTX::INT_PTX_LDU_G_v2f64_ELE_avar);
+ NVPTX::INT_PTX_LDU_G_v2i8_ELE_asi,
+ NVPTX::INT_PTX_LDU_G_v2i16_ELE_asi,
+ NVPTX::INT_PTX_LDU_G_v2i32_ELE_asi,
+ NVPTX::INT_PTX_LDU_G_v2i64_ELE_asi,
+ NVPTX::INT_PTX_LDU_G_v2f32_ELE_asi,
+ NVPTX::INT_PTX_LDU_G_v2f64_ELE_asi);
break;
case NVPTXISD::LoadV4:
Opcode = pickOpcodeForVT(
- EltVT.getSimpleVT().SimpleTy, NVPTX::INT_PTX_LDG_G_v4i8_ELE_avar,
- NVPTX::INT_PTX_LDG_G_v4i16_ELE_avar,
- NVPTX::INT_PTX_LDG_G_v4i32_ELE_avar, std::nullopt,
- NVPTX::INT_PTX_LDG_G_v4f32_ELE_avar, std::nullopt);
+ EltVT.getSimpleVT().SimpleTy, NVPTX::INT_PTX_LDG_G_v4i8_ELE_asi,
+ NVPTX::INT_PTX_LDG_G_v4i16_ELE_asi,
+ NVPTX::INT_PTX_LDG_G_v4i32_ELE_asi, std::nullopt,
+ NVPTX::INT_PTX_LDG_G_v4f32_ELE_asi, std::nullopt);
break;
case NVPTXISD::LDUV4:
Opcode = pickOpcodeForVT(
- EltVT.getSimpleVT().SimpleTy, NVPTX::INT_PTX_LDU_G_v4i8_ELE_avar,
- NVPTX::INT_PTX_LDU_G_v4i16_ELE_avar,
- NVPTX::INT_PTX_LDU_G_v4i32_ELE_avar, std::nullopt,
- NVPTX::INT_PTX_LDU_G_v4f32_ELE_avar, std::nullopt);
+ EltVT.getSimpleVT().SimpleTy, NVPTX::INT_PTX_LDU_G_v4i8_ELE_asi,
+ NVPTX::INT_PTX_LDU_G_v4i16_ELE_asi,
+ NVPTX::INT_PTX_LDU_G_v4i32_ELE_asi, std::nullopt,
+ NVPTX::INT_PTX_LDU_G_v4f32_ELE_asi, std::nullopt);
break;
}
- if (!Opcode)
- return false;
- SDValue Ops[] = { Addr, Chain };
- LD = CurDAG->getMachineNode(*Opcode, DL, InstVTList, Ops);
} else {
if (TM.is64Bit()) {
SelectADDRri64(Op1.getNode(), Op1, Base, Offset);
@@ -1369,11 +1362,11 @@ bool NVPTXDAGToDAGISel::tryLDGLDU(SDNode *N) {
break;
}
}
- if (!Opcode)
- return false;
- SDValue Ops[] = {Base, Offset, Chain};
- LD = CurDAG->getMachineNode(*Opcode, DL, InstVTList, Ops);
}
+ if (!Opcode)
+ return false;
+ SDValue Ops[] = {Base, Offset, Chain};
+ LD = CurDAG->getMachineNode(*Opcode, DL, InstVTList, Ops);
// For automatic generation of LDG (through SelectLoad[Vector], not the
// intrinsics), we may have an extending load like:
@@ -1577,7 +1570,6 @@ bool NVPTXDAGToDAGISel::tryStoreVector(SDNode *N) {
std::nullopt, NVPTX::STV_f32_v4_asi, std::nullopt);
break;
}
- Ops.append({Base, Offset});
} else {
if (PointerSize == 64) {
SelectADDRri64(N2.getNode(), N2, Base, Offset);
@@ -1617,12 +1609,10 @@ bool NVPTXDAGToDAGISel::tryStoreVector(SDNode *N) {
break;
}
}
- Ops.append({Base, Offset});
}
if (!Opcode)
return false;
-
- Ops.push_back(Chain);
+ Ops.append({Base, Offset, Chain});
ST = CurDAG->getMachineNode(*Opcode, DL, MVT::Other, Ops);
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index c98d8a3e22b43..3373f9e90524f 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -2718,23 +2718,23 @@ defm INT_PTX_SATOM_XOR : ATOM2_bitwise_impl<"xor">;
// Scalar
multiclass LDU_G<string TyStr, NVPTXRegClass regclass> {
- def avar: NVPTXInst<(outs regclass:$result), (ins imemAny:$src),
- !strconcat("ldu.global.", TyStr),
+ def asi: NVPTXInst<(outs regclass:$result), (ins imemAny:$src, Offseti32imm:$offset),
+ "ldu.global." # TyStr # " \t$result, [$src$offset];",
[]>, Requires<[hasLDU]>;
def ari : NVPTXInst<(outs regclass:$result), (ins MEMri:$src),
- !strconcat("ldu.global.", TyStr),
+ "ldu.global." # TyStr # " \t$result, [$src];",
[]>, Requires<[hasLDU]>;
def ari64 : NVPTXInst<(outs regclass:$result), (ins MEMri64:$src),
- !strconcat("ldu.global.", TyStr),
+ "ldu.global." # TyStr # " \t$result, [$src];",
[]>, Requires<[hasLDU]>;
}
-defm INT_PTX_LDU_GLOBAL_i8 : LDU_G<"u8 \t$result, [$src];", Int16Regs>;
-defm INT_PTX_LDU_GLOBAL_i16 : LDU_G<"u16 \t$result, [$src];", Int16Regs>;
-defm INT_PTX_LDU_GLOBAL_i32 : LDU_G<"u32 \t$result, [$src];", Int32Regs>;
-defm INT_PTX_LDU_GLOBAL_i64 : LDU_G<"u64 \t$result, [$src];", Int64Regs>;
-defm INT_PTX_LDU_GLOBAL_f32 : LDU_G<"f32 \t$result, [$src];", Float32Regs>;
-defm INT_PTX_LDU_GLOBAL_f64 : LDU_G<"f64 \t$result, [$src];", Float64Regs>;
+defm INT_PTX_LDU_GLOBAL_i8 : LDU_G<"u8", Int16Regs>;
+defm INT_PTX_LDU_GLOBAL_i16 : LDU_G<"u16", Int16Regs>;
+defm INT_PTX_LDU_GLOBAL_i32 : LDU_G<"u32", Int32Regs>;
+defm INT_PTX_LDU_GLOBAL_i64 : LDU_G<"u64", Int64Regs>;
+defm INT_PTX_LDU_GLOBAL_f32 : LDU_G<"f32", Float32Regs>;
+defm INT_PTX_LDU_GLOBAL_f64 : LDU_G<"f64", Float64Regs>;
// vector
@@ -2742,56 +2742,40 @@ defm INT_PTX_LDU_GLOBAL_f64 : LDU_G<"f64 \t$result, [$src];", Float64Regs>;
multiclass VLDU_G_ELE_V2<string TyStr, NVPTXRegClass regclass> {
def _ari32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
(ins MEMri:$src),
- !strconcat("ldu.global.", TyStr), []>;
+ "ldu.global.v2." # TyStr # " \t{{$dst1, $dst2}}, [$src];", []>;
def _ari64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
(ins MEMri64:$src),
- !strconcat("ldu.global.", TyStr), []>;
- def _avar: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
- (ins imemAny:$src),
- !strconcat("ldu.global.", TyStr), []>;
+ "ldu.global.v2." # TyStr # " \t{{$dst1, $dst2}}, [$src];", []>;
+ def _asi: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
+ (ins imemAny:$src, Offseti32imm:$offset),
+ "ldu.global.v2." # TyStr # " \t{{$dst1, $dst2}}, [$src$offset];", []>;
}
multiclass VLDU_G_ELE_V4<string TyStr, NVPTXRegClass regclass> {
def _ari32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
regclass:$dst4), (ins MEMri:$src),
- !strconcat("ldu.global.", TyStr), []>;
+ "ldu.global.v4." # TyStr # " \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", []>;
def _ari64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
regclass:$dst4), (ins MEMri64:$src),
- !strconcat("ldu.global.", TyStr), []>;
- def _avar: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
- regclass:$dst4), (ins imemAny:$src),
- !strconcat("ldu.global.", TyStr), []>;
-}
-
-defm INT_PTX_LDU_G_v2i8_ELE
- : VLDU_G_ELE_V2<"v2.u8 \t{{$dst1, $dst2}}, [$src];", Int16Regs>;
-defm INT_PTX_LDU_G_v2i16_ELE
- : VLDU_G_ELE_V2<"v2.u16 \t{{$dst1, $dst2}}, [$src];", Int16Regs>;
-defm INT_PTX_LDU_G_v2i32_ELE
- : VLDU_G_ELE_V2<"v2.u32 \t{{$dst1, $dst2}}, [$src];", Int32Regs>;
-defm INT_PTX_LDU_G_v2f32_ELE
- : VLDU_G_ELE_V2<"v2.f32 \t{{$dst1, $dst2}}, [$src];", Float32Regs>;
-defm INT_PTX_LDU_G_v2i64_ELE
- : VLDU_G_ELE_V2<"v2.u64 \t{{$dst1, $dst2}}, [$src];", Int64Regs>;
-defm INT_PTX_LDU_G_v2f64_ELE
- : VLDU_G_ELE_V2<"v2.f64 \t{{$dst1, $dst2}}, [$src];", Float64Regs>;
-defm INT_PTX_LDU_G_v4i8_ELE
- : VLDU_G_ELE_V4<"v4.u8 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int16Regs>;
-defm INT_PTX_LDU_G_v4i16_ELE
- : VLDU_G_ELE_V4<"v4.u16 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];",
- Int16Regs>;
-defm INT_PTX_LDU_G_v4i32_ELE
- : VLDU_G_ELE_V4<"v4.u32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];",
- Int32Regs>;
-defm INT_PTX_LDU_G_v4f16_ELE
- : VLDU_G_ELE_V4<"v4.b16 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];",
- Int16Regs>;
-defm INT_PTX_LDU_G_v4f16x2_ELE
- : VLDU_G_ELE_V4<"v4.b32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];",
- Int32Regs>;
-defm INT_PTX_LDU_G_v4f32_ELE
- : VLDU_G_ELE_V4<"v4.f32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];",
- Float32Regs>;
+ "ldu.global.v4." # TyStr # " \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", []>;
+ def _asi: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
+ regclass:$dst4), (ins imemAny:$src, Offseti32imm:$offset),
+ "ldu.global.v4." # TyStr # " \t{{$dst1, $dst2, $dst3, $dst4}}, [$src$offset];", []>;
+}
+
+defm INT_PTX_LDU_G_v2i8_ELE : VLDU_G_ELE_V2<"u8", Int16Regs>;
+defm INT_PTX_LDU_G_v2i16_ELE : VLDU_G_ELE_V2<"u16", Int16Regs>;
+defm INT_PTX_LDU_G_v2i32_ELE : VLDU_G_ELE_V2<"u32", Int32Regs>;
+defm INT_PTX_LDU_G_v2f32_ELE : VLDU_G_ELE_V2<"f32", Float32Regs>;
+defm INT_PTX_LDU_G_v2i64_ELE : VLDU_G_ELE_V2<"u64", Int64Regs>;
+defm INT_PTX_LDU_G_v2f64_ELE : VLDU_G_ELE_V2<"f64", Float64Regs>;
+
+defm INT_PTX_LDU_G_v4i8_ELE : VLDU_G_ELE_V4<"u8", Int16Regs>;
+defm INT_PTX_LDU_G_v4i16_ELE : VLDU_G_ELE_V4<"u16", Int16Regs>;
+defm INT_PTX_LDU_G_v4i32_ELE : VLDU_G_ELE_V4<"u32", Int32Regs>;
+defm INT_PTX_LDU_G_v4f16_ELE : VLDU_G_ELE_V4<"b16", Int16Regs>;
+defm INT_PTX_LDU_G_v4f16x2_ELE : VLDU_G_ELE_V4<"b32", Int32Regs>;
+defm INT_PTX_LDU_G_v4f32_ELE : VLDU_G_ELE_V4<"f32", Float32Regs>;
//-----------------------------------
@@ -2803,29 +2787,23 @@ defm INT_PTX_LDU_G_v4f32_ELE
// during the lifetime of the kernel.
multiclass LDG_G<string TyStr, NVPTXRegClass regclass> {
- def avar: NVPTXInst<(outs regclass:$result), (ins imemAny:$src),
- !strconcat("ld.global.nc.", TyStr),
+ def asi: NVPTXInst<(outs regclass:$result), (ins imemAny:$src, Offseti32imm:$offset),
+ "ld.global.nc." # TyStr # " \t$result, [$src$offset];",
[]>, Requires<[hasLDG]>;
def ari : NVPTXInst<(outs regclass:$result), (ins MEMri:$src),
- !strconcat("ld.global.nc.", TyStr),
+ "ld.global.nc." # TyStr # " \t$result, [$src];",
[]>, Requires<[hasLDG]>;
def ari64 : NVPTXInst<(outs regclass:$result), (ins MEMri64:$src),
- !strconcat("ld.global.nc.", TyStr),
+ "ld.global.nc." # TyStr # " \t$result, [$src];",
[]>, Requires<[hasLDG]>;
}
-defm INT_PTX_LDG_GLOBAL_i8
- : LDG_G<"u8 \t$result, [$src];", Int16Regs>;
-defm INT_PTX_LDG_GLOBAL_i16
- : LDG_G<"u16 \t$result, [$src];", Int16Regs>;
-defm INT_PTX_LDG_GLOBAL_i32
- : LDG_G<"u32 \t$result, [$src];", Int32Regs>;
-defm INT_PTX_LDG_GLOBAL_i64
- : LDG_G<"u64 \t$result, [$src];", Int64Regs>;
-defm INT_PTX_LDG_GLOBAL_f32
- : LDG_G<"f32 \t$result, [$src];", Float32Regs>;
-defm INT_PTX_LDG_GLOBAL_f64
- : LDG_G<"f64 \t$result, [$src];", Float64Regs>;
+defm INT_PTX_LDG_GLOBAL_i8 : LDG_G<"u8", Int16Regs>;
+defm INT_PTX_LDG_GLOBAL_i16 : LDG_G<"u16", Int16Regs>;
+defm INT_PTX_LDG_GLOBAL_i32 : LDG_G<"u32", Int32Regs>;
+defm INT_PTX_LDG_GLOBAL_i64 : LDG_G<"u64", Int64Regs>;
+defm INT_PTX_LDG_GLOBAL_f32 : LDG_G<"f32", Float32Regs>;
+defm INT_PTX_LDG_GLOBAL_f64 : LDG_G<"f64", Float64Regs>;
// vector
@@ -2833,54 +2811,39 @@ defm INT_PTX_LDG_GLOBAL_f64
multiclass VLDG_G_ELE_V2<string TyStr, NVPTXRegClass regclass> {
def _ari32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
(ins MEMri:$src),
- !strconcat("ld.global.nc.", TyStr), []>;
+ "ld.global.nc.v2." # TyStr # " \t{{$dst1, $dst2}}, [$src];", []>;
def _ari64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
(ins MEMri64:$src),
- !strconcat("ld.global.nc.", TyStr), []>;
- def _avar: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
- (ins imemAny:$src),
- !strconcat("ld.global.nc.", TyStr), []>;
+ "ld.global.nc.v2." # TyStr # " \t{{$dst1, $dst2}}, [$src];", []>;
+ def _asi: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
+ (ins imemAny:$src, Offseti32imm:$offset),
+ "ld.global.nc.v2." # TyStr # " \t{{$dst1, $dst2}}, [$src$offset];", []>;
}
multiclass VLDG_G_ELE_V4<string TyStr, NVPTXRegClass regclass> {
- def _areg32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
- regclass:$dst4), (ins Int32Regs:$src),
- !strconcat("ld.global.nc.", TyStr), []>;
- def _areg64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
- regclass:$dst4), (ins Int64Regs:$src),
- !strconcat("ld.global.nc.", TyStr), []>;
def _ari32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
regclass:$dst4), (ins MEMri:$src),
- !strconcat("ld.global.nc.", TyStr), []>;
+ "ld.global.nc.v4." # TyStr # " \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", []>;
def _ari64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
regclass:$dst4), (ins MEMri64:$src),
- !strconcat("ld.global.nc.", TyStr), []>;
- def _avar: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
- regclass:$dst4), (ins imemAny:$src),
- !strconcat("ld.global.nc.", TyStr), []>;
+ "ld.global.nc.v4." # TyStr # " \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", []>;
+ def _asi: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
+ regclass:$dst4), (ins imemAny:$src, Offseti32imm:$offset),
+ "ld.global.nc.v4." # TyStr # " \t{{$dst1, $dst2, $dst3, $dst4}}, [$src$offset];", []>;
}
// FIXME: 8-bit LDG should be fixed once LDG/LDU nodes are made into proper loads.
-defm INT_PTX_LDG_G_v2i8_ELE
- : VLDG_G_ELE_V2<"v2.u8 \t{{$dst1, $dst2}}, [$src];", Int16Regs>;
-defm INT_PTX_LDG_G_v2i16_ELE
- : VLDG_G_ELE_V2<"v2.u16 \t{{$dst1, $dst2}}, [$src];", Int16Regs>;
-defm INT_PTX_LDG_G_v2i32_ELE
- : VLDG_G_ELE_V2<"v2.u32 \t{{$dst1, $dst2}}, [$src];", Int32Regs>;
-defm INT_PTX_LDG_G_v2f32_ELE
- : VLDG_G_ELE_V2<"v2.f32 \t{{$dst1, $dst2}}, [$src];", Float32Regs>;
-defm INT_PTX_LDG_G_v2i64_ELE
- : VLDG_G_ELE_V2<"v2.u64 \t{{$dst1, $dst2}}, [$src];", Int64Regs>;
-defm INT_PTX_LDG_G_v2f64_ELE
- : VLDG_G_ELE_V2<"v2.f64 \t{{$dst1, $dst2}}, [$src];", Float64Regs>;
-defm INT_PTX_LDG_G_v4i8_ELE
- : VLDG_G_ELE_V4<"v4.u8 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int16Regs>;
-defm INT_PTX_LDG_G_v4i16_ELE
- : VLDG_G_ELE_V4<"v4.u16 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int16Regs>;
-defm INT_PTX_LDG_G_v4i32_ELE
- : VLDG_G_ELE_V4<"v4.u32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int32Regs>;
-defm INT_PTX_LDG_G_v4f32_ELE
- : VLDG_G_ELE_V4<"v4.f32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Float32Regs>;
+defm INT_PTX_LDG_G_v2i8_ELE : VLDG_G_ELE_V2<"u8", Int16Regs>;
+defm INT_PTX_LDG_G_v2i16_ELE : VLDG_G_ELE_V2<"u16", Int16Regs>;
+defm INT_PTX_LDG_G_v2i32_ELE : VLDG_G_ELE_V2<"u32", Int32Regs>;
+defm INT_PTX_LDG_G_v2f32_ELE : VLDG_G_ELE_V2<"f32", Float32Regs>;
+defm INT_PTX_LDG_G_v2i64_ELE : VLDG_G_ELE_V2<"u64", Int64Regs>;
+defm INT_PTX_LDG_G_v2f64_ELE : VLDG_G_ELE_V2<"f64", Float64Regs>;
+
+defm INT_PTX_LDG_G_v4i8_ELE : VLDG_G_ELE_V4<"u8", Int16Regs>;
+defm INT_PTX_LDG_G_v4i16_ELE : VLDG_G_ELE_V4<"u16", Int16Regs>;
+defm INT_PTX_LDG_G_v4i32_ELE : VLDG_G_ELE_V4<"u32", Int32Regs>;
+defm INT_PTX_LDG_G_v4f32_ELE : VLDG_G_ELE_V4<"f32", Float32Regs>;
multiclass NG_TO_G<string Str> {
diff --git a/llvm/test/CodeGen/NVPTX/ldu-ldg.ll b/llvm/test/CodeGen/NVPTX/ldu-ldg.ll
index c144de412c0cb..4c5c44a9bf44d 100644
--- a/llvm/test/CodeGen/NVPTX/ldu-ldg.ll
+++ b/llvm/test/CodeGen/NVPTX/ldu-ldg.ll
@@ -1,3 +1,4 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_32 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_32 | %ptxas-verify %}
@@ -22,129 +23,305 @@ declare double @llvm.nvvm.ldg.global.f.f64.p1(ptr addrspace(1) %ptr, i32 %align)
declare half @llvm.nvvm.ldg.global.f.f16.p1(ptr addrspace(1) %ptr, i32 %align)
declare <2 x half> @llvm.nvvm.ldg.global.f.v2f16.p1(ptr addrspace(1) %ptr, i32 %align)
-; CHECK-LABEL: test_ldu_i8
define i8 @test_ldu_i8(ptr addrspace(1) %ptr) {
- ; CHECK: ldu.global.u8
+; CHECK-LABEL: test_ldu_i8(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u64 %rd1, [test_ldu_i8_param_0];
+; CHECK-NEXT: ldu.global.u8 %rs1, [%rd1];
+; CHECK-NEXT: cvt.u32.u16 %r1, %rs1;
+; CHECK-NEXT: and.b32 %r2, %r1, 255;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
%val = tail call i8 @llvm.nvvm.ldu.global.i.i8.p1(ptr addrspace(1) %ptr, i32 4)
ret i8 %val
}
-; CHECK-LABEL: test_ldu_i16
define i16 @test_ldu_i16(ptr addrspace(1) %ptr) {
- ; CHECK: ldu.global.u16
+; CHECK-LABEL: test_ldu_i16(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u64 %rd1, [test_ldu_i16_param_0];
+; CHECK-NEXT: ldu.global.u16 %rs1, [%rd1];
+; CHECK-NEXT: cvt.u32.u16 %r1, %rs1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
%val = tail call i16 @llvm.nvvm.ldu.global.i.i16.p1(ptr addrspace(1) %ptr, i32 2)
ret i16 %val
}
-; CHECK-LABEL: test_ldu_i32
define i32 @test_ldu_i32(ptr addrspace(1) %ptr) {
- ; CHECK: ldu.global.u32
+; CHECK-LABEL: test_ldu_i32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u64 %rd1, [test_ldu_i32_param_0];
+; CHECK-NEXT: ldu.global.u32 %r1, [%rd1];
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
%val = tail call i32 @llvm.nvvm.ldu.global.i.i32.p1(ptr addrspace(1) %ptr, i32 4)
ret i32 %val
}
-; CHECK-LABEL: test_ldu_i64
define i64 @test_ldu_i64(ptr addrspace(1) %ptr) {
- ; CHECK: ldu.global.u64
+; CHECK-LABEL: test_ldu_i64(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u64 %rd1, [test_ldu_i64_param_0];
+; CHECK-NEXT: ldu.global.u64 %rd2, [%rd1];
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
+; CHECK-NEXT: ret;
%val = tail call i64 @llvm.nvvm.ldu.global.i.i64.p1(ptr addrspace(1) %ptr, i32 8)
ret i64 %val
}
-; CHECK-LABEL: test_ldu_p
define ptr @test_ldu_p(ptr addrspace(1) %ptr) {
- ; CHECK: ldu.global.u64
+; CHECK-LABEL: test_ldu_p(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u64 %rd1, [test_ldu_p_param_0];
+; CHECK-NEXT: ldu.global.u64 %rd2, [%rd1];
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
+; CHECK-NEXT: ret;
%val = tail call ptr @llvm.nvvm.ldu.global.p.p1(ptr addrspace(1) %ptr, i32 8)
ret ptr %val
}
-
-; CHECK-LABEL: test_ldu_f32
define float @test_ldu_f32(ptr addrspace(1) %ptr) {
- ; CHECK: ldu.global.f32
+; CHECK-LABEL: test_ldu_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u64 %rd1, [test_ldu_f32_param_0];
+; CHECK-NEXT: ldu.global.f32 %f1, [%rd1];
+; CHECK-NEXT: st.param.f32 [func_retval0], %f1;
+; CHECK-NEXT: ret;
%val = tail call float @llvm.nvvm.ldu.global.f.f32.p1(ptr addrspace(1) %ptr, i32 4)
ret float %val
}
-; CHECK-LABEL: test_ldu_f64
define double @test_ldu_f64(ptr addrspace(1) %ptr) {
- ; CHECK: ldu.global.f64
+; CHECK-LABEL: test_ldu_f64(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-NEXT: .reg .f64 %fd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u64 %rd1, [test_ldu_f64_param_0];
+; CHECK-NEXT: ldu.global.f64 %fd1, [%rd1];
+; CHECK-NEXT: st.param.f64 [func_retval0], %fd1;
+; CHECK-NEXT: ret;
%val = tail call double @llvm.nvvm.ldu.global.f.f64.p1(ptr addrspace(1) %ptr, i32 8)
ret double %val
}
-; CHECK-LABEL: test_ldu_f16
define half @test_ldu_f16(ptr addrspace(1) %ptr) {
- ; CHECK: ldu.global.u16
+; CHECK-LABEL: test_ldu_f16(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u64 %rd1, [test_ldu_f16_param_0];
+; CHECK-NEXT: ldu.global.u16 %rs1, [%rd1];
+; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;
+; CHECK-NEXT: ret;
%val = tail call half @llvm.nvvm.ldu.global.f.f16.p1(ptr addrspace(1) %ptr, i32 2)
ret half %val
}
-; CHECK-LABEL: test_ldu_v2f16
define <2 x half> @test_ldu_v2f16(ptr addrspace(1) %ptr) {
- ; CHECK: ldu.global.u32
+; CHECK-LABEL: test_ldu_v2f16(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u64 %rd1, [test_ldu_v2f16_param_0];
+; CHECK-NEXT: ldu.global.u32 %r1, [%rd1];
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
%val = tail call <2 x half> @llvm.nvvm.ldu.global.f.v2f16.p1(ptr addrspace(1) %ptr, i32 4)
ret <2 x half> %val
}
-; CHECK-LABEL: test_ldg_i8
define i8 @test_ldg_i8(ptr addrspace(1) %ptr) {
- ; CHECK: ld.global.nc.u8
+; CHECK-LABEL: test_ldg_i8(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u64 %rd1, [test_ldg_i8_param_0];
+; CHECK-NEXT: ld.global.nc.u8 %rs1, [%rd1];
+; CHECK-NEXT: cvt.u32.u8 %r1, %rs1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
%val = tail call i8 @llvm.nvvm.ldg.global.i.i8.p1(ptr addrspace(1) %ptr, i32 4)
ret i8 %val
}
-; CHECK-LABEL: test_ldg_i16
define i16 @test_ldg_i16(ptr addrspace(1) %ptr) {
- ; CHECK: ld.global.nc.u16
+; CHECK-LABEL: test_ldg_i16(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u64 %rd1, [test_ldg_i16_param_0];
+; CHECK-NEXT: ld.global.nc.u16 %rs1, [%rd1];
+; CHECK-NEXT: cvt.u32.u16 %r1, %rs1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
%val = tail call i16 @llvm.nvvm.ldg.global.i.i16.p1(ptr addrspace(1) %ptr, i32 2)
ret i16 %val
}
-; CHECK-LABEL: test_ldg_i32
define i32 @test_ldg_i32(ptr addrspace(1) %ptr) {
- ; CHECK: ld.global.nc.u32
+; CHECK-LABEL: test_ldg_i32(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u64 %rd1, [test_ldg_i32_param_0];
+; CHECK-NEXT: ld.global.nc.u32 %r1, [%rd1];
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
%val = tail call i32 @llvm.nvvm.ldg.global.i.i32.p1(ptr addrspace(1) %ptr, i32 4)
ret i32 %val
}
-; CHECK-LABEL: test_ldg_i64
define i64 @test_ldg_i64(ptr addrspace(1) %ptr) {
- ; CHECK: ld.global.nc.u64
+; CHECK-LABEL: test_ldg_i64(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u64 %rd1, [test_ldg_i64_param_0];
+; CHECK-NEXT: ld.global.nc.u64 %rd2, [%rd1];
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
+; CHECK-NEXT: ret;
%val = tail call i64 @llvm.nvvm.ldg.global.i.i64.p1(ptr addrspace(1) %ptr, i32 8)
ret i64 %val
}
-; CHECK-LABEL: test_ldg_p
define ptr @test_ldg_p(ptr addrspace(1) %ptr) {
- ; CHECK: ld.global.nc.u64
+; CHECK-LABEL: test_ldg_p(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u64 %rd1, [test_ldg_p_param_0];
+; CHECK-NEXT: ld.global.nc.u64 %rd2, [%rd1];
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd2;
+; CHECK-NEXT: ret;
%val = tail call ptr @llvm.nvvm.ldg.global.p.p1(ptr addrspace(1) %ptr, i32 8)
ret ptr %val
}
-; CHECK-LABEL: test_ldg_f32
define float @test_ldg_f32(ptr addrspace(1) %ptr) {
- ; CHECK: ld.global.nc.f32
+; CHECK-LABEL: test_ldg_f32(
+; CHECK: {
+; CHECK-NEXT: .reg .f32 %f<2>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u64 %rd1, [test_ldg_f32_param_0];
+; CHECK-NEXT: ld.global.nc.f32 %f1, [%rd1];
+; CHECK-NEXT: st.param.f32 [func_retval0], %f1;
+; CHECK-NEXT: ret;
%val = tail call float @llvm.nvvm.ldg.global.f.f32.p1(ptr addrspace(1) %ptr, i32 4)
ret float %val
}
-; CHECK-LABEL: test_ldg_f64
define double @test_ldg_f64(ptr addrspace(1) %ptr) {
- ; CHECK: ld.global.nc.f64
+; CHECK-LABEL: test_ldg_f64(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-NEXT: .reg .f64 %fd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u64 %rd1, [test_ldg_f64_param_0];
+; CHECK-NEXT: ld.global.nc.f64 %fd1, [%rd1];
+; CHECK-NEXT: st.param.f64 [func_retval0], %fd1;
+; CHECK-NEXT: ret;
%val = tail call double @llvm.nvvm.ldg.global.f.f64.p1(ptr addrspace(1) %ptr, i32 8)
ret double %val
}
-; CHECK-LABEL: test_ldg_f16
define half @test_ldg_f16(ptr addrspace(1) %ptr) {
- ; CHECK: ld.global.nc.u16
+; CHECK-LABEL: test_ldg_f16(
+; CHECK: {
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u64 %rd1, [test_ldg_f16_param_0];
+; CHECK-NEXT: ld.global.nc.u16 %rs1, [%rd1];
+; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;
+; CHECK-NEXT: ret;
%val = tail call half @llvm.nvvm.ldg.global.f.f16.p1(ptr addrspace(1) %ptr, i32 2)
ret half %val
}
-; CHECK-LABEL: test_ldg_v2f16
define <2 x half> @test_ldg_v2f16(ptr addrspace(1) %ptr) {
- ; CHECK: ld.global.nc.u32
+; CHECK-LABEL: test_ldg_v2f16(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u64 %rd1, [test_ldg_v2f16_param_0];
+; CHECK-NEXT: ld.global.nc.u32 %r1, [%rd1];
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
%val = tail call <2 x half> @llvm.nvvm.ldg.global.f.v2f16.p1(ptr addrspace(1) %ptr, i32 4)
ret <2 x half> %val
}
+
+ at g = addrspace(1) global i32 0
+
+define i32 @test_ldg_asi() {
+; CHECK-LABEL: test_ldg_asi(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.global.nc.u32 %r1, [g+4];
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %val = tail call i32 @llvm.nvvm.ldg.global.i.i32.p1(ptr addrspace(1) getelementptr (i8, ptr addrspace(1) @g, i32 4), i32 4)
+ ret i32 %val
+}
+
+define i32 @test_lug_asi() {
+; CHECK-LABEL: test_lug_asi(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ldu.global.u32 %r1, [g+4];
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT: ret;
+ %val = tail call i32 @llvm.nvvm.ldu.global.i.i32.p1(ptr addrspace(1) getelementptr (i8, ptr addrspace(1) @g, i32 4), i32 4)
+ ret i32 %val
+}
diff --git a/llvm/test/CodeGen/NVPTX/variadics-backend.ll b/llvm/test/CodeGen/NVPTX/variadics-backend.ll
index 4d4db21c6ed0d..377528b94f505 100644
--- a/llvm/test/CodeGen/NVPTX/variadics-backend.ll
+++ b/llvm/test/CodeGen/NVPTX/variadics-backend.ll
@@ -214,34 +214,33 @@ define dso_local i32 @bar() {
; CHECK-PTX-NEXT: .reg .b64 %SPL;
; CHECK-PTX-NEXT: .reg .b16 %rs<8>;
; CHECK-PTX-NEXT: .reg .b32 %r<4>;
-; CHECK-PTX-NEXT: .reg .b64 %rd<6>;
+; CHECK-PTX-NEXT: .reg .b64 %rd<5>;
; CHECK-PTX-EMPTY:
; CHECK-PTX-NEXT: // %bb.0: // %entry
; CHECK-PTX-NEXT: mov.u64 %SPL, __local_depot3;
; CHECK-PTX-NEXT: cvta.local.u64 %SP, %SPL;
; CHECK-PTX-NEXT: add.u64 %rd2, %SPL, 0;
-; CHECK-PTX-NEXT: mov.u64 %rd3, __const_$_bar_$_s1;
-; CHECK-PTX-NEXT: ld.global.nc.u8 %rs1, [%rd3+7];
+; CHECK-PTX-NEXT: ld.global.nc.u8 %rs1, [__const_$_bar_$_s1+7];
; CHECK-PTX-NEXT: cvt.u16.u8 %rs2, %rs1;
; CHECK-PTX-NEXT: st.local.u8 [%rd2+2], %rs2;
-; CHECK-PTX-NEXT: ld.global.nc.u8 %rs3, [%rd3+6];
+; CHECK-PTX-NEXT: ld.global.nc.u8 %rs3, [__const_$_bar_$_s1+6];
; CHECK-PTX-NEXT: cvt.u16.u8 %rs4, %rs3;
; CHECK-PTX-NEXT: st.local.u8 [%rd2+1], %rs4;
-; CHECK-PTX-NEXT: ld.global.nc.u8 %rs5, [%rd3+5];
+; CHECK-PTX-NEXT: ld.global.nc.u8 %rs5, [__const_$_bar_$_s1+5];
; CHECK-PTX-NEXT: cvt.u16.u8 %rs6, %rs5;
; CHECK-PTX-NEXT: st.local.u8 [%rd2], %rs6;
; CHECK-PTX-NEXT: mov.b32 %r1, 1;
; CHECK-PTX-NEXT: st.u32 [%SP+8], %r1;
; CHECK-PTX-NEXT: mov.b16 %rs7, 1;
; CHECK-PTX-NEXT: st.u8 [%SP+12], %rs7;
-; CHECK-PTX-NEXT: mov.b64 %rd4, 1;
-; CHECK-PTX-NEXT: st.u64 [%SP+16], %rd4;
-; CHECK-PTX-NEXT: add.u64 %rd5, %SP, 8;
+; CHECK-PTX-NEXT: mov.b64 %rd3, 1;
+; CHECK-PTX-NEXT: st.u64 [%SP+16], %rd3;
+; CHECK-PTX-NEXT: add.u64 %rd4, %SP, 8;
; CHECK-PTX-NEXT: { // callseq 1, 0
; CHECK-PTX-NEXT: .param .b32 param0;
; CHECK-PTX-NEXT: st.param.b32 [param0], 1;
; CHECK-PTX-NEXT: .param .b64 param1;
-; CHECK-PTX-NEXT: st.param.b64 [param1], %rd5;
+; CHECK-PTX-NEXT: st.param.b64 [param1], %rd4;
; CHECK-PTX-NEXT: .param .b32 retval0;
; CHECK-PTX-NEXT: call.uni (retval0),
; CHECK-PTX-NEXT: variadics2,
@@ -380,28 +379,27 @@ define dso_local void @qux() {
; CHECK-PTX-NEXT: .reg .b64 %SP;
; CHECK-PTX-NEXT: .reg .b64 %SPL;
; CHECK-PTX-NEXT: .reg .b32 %r<3>;
-; CHECK-PTX-NEXT: .reg .b64 %rd<10>;
+; CHECK-PTX-NEXT: .reg .b64 %rd<9>;
; CHECK-PTX-EMPTY:
; CHECK-PTX-NEXT: // %bb.0: // %entry
; CHECK-PTX-NEXT: mov.u64 %SPL, __local_depot7;
; CHECK-PTX-NEXT: cvta.local.u64 %SP, %SPL;
; CHECK-PTX-NEXT: add.u64 %rd2, %SPL, 0;
-; CHECK-PTX-NEXT: mov.u64 %rd3, __const_$_qux_$_s;
-; CHECK-PTX-NEXT: ld.global.nc.u64 %rd4, [%rd3+8];
-; CHECK-PTX-NEXT: st.local.u64 [%rd2+8], %rd4;
-; CHECK-PTX-NEXT: ld.global.nc.u64 %rd5, [__const_$_qux_$_s];
-; CHECK-PTX-NEXT: st.local.u64 [%rd2], %rd5;
-; CHECK-PTX-NEXT: mov.b64 %rd6, 1;
-; CHECK-PTX-NEXT: st.u64 [%SP+16], %rd6;
-; CHECK-PTX-NEXT: ld.local.u64 %rd7, [%rd2];
-; CHECK-PTX-NEXT: ld.local.u64 %rd8, [%rd2+8];
-; CHECK-PTX-NEXT: add.u64 %rd9, %SP, 16;
+; CHECK-PTX-NEXT: ld.global.nc.u64 %rd3, [__const_$_qux_$_s+8];
+; CHECK-PTX-NEXT: st.local.u64 [%rd2+8], %rd3;
+; CHECK-PTX-NEXT: ld.global.nc.u64 %rd4, [__const_$_qux_$_s];
+; CHECK-PTX-NEXT: st.local.u64 [%rd2], %rd4;
+; CHECK-PTX-NEXT: mov.b64 %rd5, 1;
+; CHECK-PTX-NEXT: st.u64 [%SP+16], %rd5;
+; CHECK-PTX-NEXT: ld.local.u64 %rd6, [%rd2];
+; CHECK-PTX-NEXT: ld.local.u64 %rd7, [%rd2+8];
+; CHECK-PTX-NEXT: add.u64 %rd8, %SP, 16;
; CHECK-PTX-NEXT: { // callseq 3, 0
; CHECK-PTX-NEXT: .param .align 8 .b8 param0[16];
-; CHECK-PTX-NEXT: st.param.b64 [param0], %rd7;
-; CHECK-PTX-NEXT: st.param.b64 [param0+8], %rd8;
+; CHECK-PTX-NEXT: st.param.b64 [param0], %rd6;
+; CHECK-PTX-NEXT: st.param.b64 [param0+8], %rd7;
; CHECK-PTX-NEXT: .param .b64 param1;
-; CHECK-PTX-NEXT: st.param.b64 [param1], %rd9;
+; CHECK-PTX-NEXT: st.param.b64 [param1], %rd8;
; CHECK-PTX-NEXT: .param .b32 retval0;
; CHECK-PTX-NEXT: call.uni (retval0),
; CHECK-PTX-NEXT: variadics4,
More information about the llvm-commits
mailing list