[llvm] r331941 - [NVPTX] Added a feature to use short pointers for const/local/shared AS.
Artem Belevich via llvm-commits
llvm-commits at lists.llvm.org
Wed May 9 16:46:19 PDT 2018
Author: tra
Date: Wed May 9 16:46:19 2018
New Revision: 331941
URL: http://llvm.org/viewvc/llvm-project?rev=331941&view=rev
Log:
[NVPTX] Added a feature to use short pointers for const/local/shared AS.
Const/local/shared address spaces are all < 4GB and we can always use
32-bit pointers to access them. This has substantial performance impact
on kernels that uses shared memory for intermediary results.
The feature is disabled by default.
Differential Revision: https://reviews.llvm.org/D46147
Modified:
llvm/trunk/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
llvm/trunk/lib/Target/NVPTX/NVPTXISelDAGToDAG.h
llvm/trunk/lib/Target/NVPTX/NVPTXISelLowering.cpp
llvm/trunk/lib/Target/NVPTX/NVPTXInstrInfo.td
llvm/trunk/lib/Target/NVPTX/NVPTXIntrinsics.td
llvm/trunk/lib/Target/NVPTX/NVPTXSubtarget.h
llvm/trunk/lib/Target/NVPTX/NVPTXTargetMachine.cpp
llvm/trunk/lib/Target/NVPTX/NVPTXTargetMachine.h
llvm/trunk/test/CodeGen/NVPTX/addrspacecast.ll
llvm/trunk/test/CodeGen/NVPTX/ld-addrspace.ll
llvm/trunk/test/CodeGen/NVPTX/st-addrspace.ll
Modified: llvm/trunk/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp?rev=331941&r1=331940&r2=331941&view=diff
==============================================================================
--- llvm/trunk/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp (original)
+++ llvm/trunk/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp Wed May 9 16:46:19 2018
@@ -66,6 +66,10 @@ bool NVPTXDAGToDAGISel::allowUnsafeFPMat
return TL->allowUnsafeFPMath(*MF);
}
+bool NVPTXDAGToDAGISel::useShortPointers() const {
+ return TM.useShortPointers();
+}
+
/// Select - Select instructions not customized! Used for
/// expanded, promoted and normal instructions.
void NVPTXDAGToDAGISel::Select(SDNode *N) {
@@ -732,7 +736,6 @@ void NVPTXDAGToDAGISel::SelectAddrSpaceC
AddrSpaceCastSDNode *CastN = cast<AddrSpaceCastSDNode>(N);
unsigned SrcAddrSpace = CastN->getSrcAddressSpace();
unsigned DstAddrSpace = CastN->getDestAddressSpace();
-
assert(SrcAddrSpace != DstAddrSpace &&
"addrspacecast must be between different address spaces");
@@ -745,13 +748,19 @@ void NVPTXDAGToDAGISel::SelectAddrSpaceC
Opc = TM.is64Bit() ? NVPTX::cvta_global_yes_64 : NVPTX::cvta_global_yes;
break;
case ADDRESS_SPACE_SHARED:
- Opc = TM.is64Bit() ? NVPTX::cvta_shared_yes_64 : NVPTX::cvta_shared_yes;
+ Opc = TM.is64Bit() ? (useShortPointers() ? NVPTX::cvta_shared_yes_6432
+ : NVPTX::cvta_shared_yes_64)
+ : NVPTX::cvta_shared_yes;
break;
case ADDRESS_SPACE_CONST:
- Opc = TM.is64Bit() ? NVPTX::cvta_const_yes_64 : NVPTX::cvta_const_yes;
+ Opc = TM.is64Bit() ? (useShortPointers() ? NVPTX::cvta_const_yes_6432
+ : NVPTX::cvta_const_yes_64)
+ : NVPTX::cvta_const_yes;
break;
case ADDRESS_SPACE_LOCAL:
- Opc = TM.is64Bit() ? NVPTX::cvta_local_yes_64 : NVPTX::cvta_local_yes;
+ Opc = TM.is64Bit() ? (useShortPointers() ? NVPTX::cvta_local_yes_6432
+ : NVPTX::cvta_local_yes_64)
+ : NVPTX::cvta_local_yes;
break;
}
ReplaceNode(N, CurDAG->getMachineNode(Opc, SDLoc(N), N->getValueType(0),
@@ -769,16 +778,19 @@ void NVPTXDAGToDAGISel::SelectAddrSpaceC
: NVPTX::cvta_to_global_yes;
break;
case ADDRESS_SPACE_SHARED:
- Opc = TM.is64Bit() ? NVPTX::cvta_to_shared_yes_64
+ Opc = TM.is64Bit() ? (useShortPointers() ? NVPTX::cvta_to_shared_yes_3264
+ : NVPTX::cvta_to_shared_yes_64)
: NVPTX::cvta_to_shared_yes;
break;
case ADDRESS_SPACE_CONST:
- Opc =
- TM.is64Bit() ? NVPTX::cvta_to_const_yes_64 : NVPTX::cvta_to_const_yes;
+ Opc = TM.is64Bit() ? (useShortPointers() ? NVPTX::cvta_to_const_yes_3264
+ : NVPTX::cvta_to_const_yes_64)
+ : NVPTX::cvta_to_const_yes;
break;
case ADDRESS_SPACE_LOCAL:
- Opc =
- TM.is64Bit() ? NVPTX::cvta_to_local_yes_64 : NVPTX::cvta_to_local_yes;
+ Opc = TM.is64Bit() ? (useShortPointers() ? NVPTX::cvta_to_local_yes_3264
+ : NVPTX::cvta_to_local_yes_64)
+ : NVPTX::cvta_to_local_yes;
break;
case ADDRESS_SPACE_PARAM:
Opc = TM.is64Bit() ? NVPTX::nvvm_ptr_gen_to_param_64
@@ -834,18 +846,20 @@ bool NVPTXDAGToDAGISel::tryLoad(SDNode *
return false;
// Address Space Setting
- unsigned int codeAddrSpace = getCodeAddrSpace(LD);
-
- if (canLowerToLDG(LD, *Subtarget, codeAddrSpace, MF)) {
+ unsigned int CodeAddrSpace = getCodeAddrSpace(LD);
+ if (canLowerToLDG(LD, *Subtarget, CodeAddrSpace, MF)) {
return tryLDGLDU(N);
}
+ unsigned int PointerSize =
+ CurDAG->getDataLayout().getPointerSizeInBits(LD->getAddressSpace());
+
// Volatile Setting
// - .volatile is only availalble for .global and .shared
bool isVolatile = LD->isVolatile();
- if (codeAddrSpace != NVPTX::PTXLdStInstCode::GLOBAL &&
- codeAddrSpace != NVPTX::PTXLdStInstCode::SHARED &&
- codeAddrSpace != NVPTX::PTXLdStInstCode::GENERIC)
+ if (CodeAddrSpace != NVPTX::PTXLdStInstCode::GLOBAL &&
+ CodeAddrSpace != NVPTX::PTXLdStInstCode::SHARED &&
+ CodeAddrSpace != NVPTX::PTXLdStInstCode::GENERIC)
isVolatile = false;
// Type Setting: fromType + fromTypeWidth
@@ -892,27 +906,27 @@ bool NVPTXDAGToDAGISel::tryLoad(SDNode *
NVPTX::LD_f32_avar, NVPTX::LD_f64_avar);
if (!Opcode)
return false;
- SDValue Ops[] = { getI32Imm(isVolatile, dl), getI32Imm(codeAddrSpace, dl),
+ SDValue Ops[] = { getI32Imm(isVolatile, dl), getI32Imm(CodeAddrSpace, dl),
getI32Imm(vecType, dl), getI32Imm(fromType, dl),
getI32Imm(fromTypeWidth, dl), Addr, Chain };
NVPTXLD = CurDAG->getMachineNode(Opcode.getValue(), dl, TargetVT,
MVT::Other, Ops);
- } else if (TM.is64Bit() ? SelectADDRsi64(N1.getNode(), N1, Base, Offset)
- : SelectADDRsi(N1.getNode(), N1, Base, Offset)) {
+ } else if (PointerSize == 64 ? SelectADDRsi64(N1.getNode(), N1, Base, Offset)
+ : SelectADDRsi(N1.getNode(), N1, Base, Offset)) {
Opcode = pickOpcodeForVT(TargetVT, NVPTX::LD_i8_asi, NVPTX::LD_i16_asi,
NVPTX::LD_i32_asi, NVPTX::LD_i64_asi,
NVPTX::LD_f16_asi, NVPTX::LD_f16x2_asi,
NVPTX::LD_f32_asi, NVPTX::LD_f64_asi);
if (!Opcode)
return false;
- SDValue Ops[] = { getI32Imm(isVolatile, dl), getI32Imm(codeAddrSpace, dl),
+ SDValue Ops[] = { getI32Imm(isVolatile, dl), getI32Imm(CodeAddrSpace, dl),
getI32Imm(vecType, dl), getI32Imm(fromType, dl),
getI32Imm(fromTypeWidth, dl), Base, Offset, Chain };
NVPTXLD = CurDAG->getMachineNode(Opcode.getValue(), dl, TargetVT,
MVT::Other, Ops);
- } else if (TM.is64Bit() ? SelectADDRri64(N1.getNode(), N1, Base, Offset)
- : SelectADDRri(N1.getNode(), N1, Base, Offset)) {
- if (TM.is64Bit())
+ } else if (PointerSize == 64 ? SelectADDRri64(N1.getNode(), N1, Base, Offset)
+ : SelectADDRri(N1.getNode(), N1, Base, Offset)) {
+ if (PointerSize == 64)
Opcode = pickOpcodeForVT(
TargetVT, NVPTX::LD_i8_ari_64, NVPTX::LD_i16_ari_64,
NVPTX::LD_i32_ari_64, NVPTX::LD_i64_ari_64, NVPTX::LD_f16_ari_64,
@@ -924,13 +938,13 @@ bool NVPTXDAGToDAGISel::tryLoad(SDNode *
NVPTX::LD_f32_ari, NVPTX::LD_f64_ari);
if (!Opcode)
return false;
- SDValue Ops[] = { getI32Imm(isVolatile, dl), getI32Imm(codeAddrSpace, dl),
+ SDValue Ops[] = { getI32Imm(isVolatile, dl), getI32Imm(CodeAddrSpace, dl),
getI32Imm(vecType, dl), getI32Imm(fromType, dl),
getI32Imm(fromTypeWidth, dl), Base, Offset, Chain };
NVPTXLD = CurDAG->getMachineNode(Opcode.getValue(), dl, TargetVT,
MVT::Other, Ops);
} else {
- if (TM.is64Bit())
+ if (PointerSize == 64)
Opcode = pickOpcodeForVT(
TargetVT, NVPTX::LD_i8_areg_64, NVPTX::LD_i16_areg_64,
NVPTX::LD_i32_areg_64, NVPTX::LD_i64_areg_64, NVPTX::LD_f16_areg_64,
@@ -943,7 +957,7 @@ bool NVPTXDAGToDAGISel::tryLoad(SDNode *
NVPTX::LD_f32_areg, NVPTX::LD_f64_areg);
if (!Opcode)
return false;
- SDValue Ops[] = { getI32Imm(isVolatile, dl), getI32Imm(codeAddrSpace, dl),
+ SDValue Ops[] = { getI32Imm(isVolatile, dl), getI32Imm(CodeAddrSpace, dl),
getI32Imm(vecType, dl), getI32Imm(fromType, dl),
getI32Imm(fromTypeWidth, dl), N1, Chain };
NVPTXLD = CurDAG->getMachineNode(Opcode.getValue(), dl, TargetVT,
@@ -977,11 +991,13 @@ bool NVPTXDAGToDAGISel::tryLoadVector(SD
// Address Space Setting
unsigned int CodeAddrSpace = getCodeAddrSpace(MemSD);
-
if (canLowerToLDG(MemSD, *Subtarget, CodeAddrSpace, MF)) {
return tryLDGLDU(N);
}
+ unsigned int PointerSize =
+ CurDAG->getDataLayout().getPointerSizeInBits(MemSD->getAddressSpace());
+
// Volatile Setting
// - .volatile is only availalble for .global and .shared
bool IsVolatile = MemSD->isVolatile();
@@ -1064,8 +1080,9 @@ bool NVPTXDAGToDAGISel::tryLoadVector(SD
getI32Imm(VecType, DL), getI32Imm(FromType, DL),
getI32Imm(FromTypeWidth, DL), Addr, Chain };
LD = CurDAG->getMachineNode(Opcode.getValue(), DL, N->getVTList(), Ops);
- } else if (TM.is64Bit() ? SelectADDRsi64(Op1.getNode(), Op1, Base, Offset)
- : SelectADDRsi(Op1.getNode(), Op1, Base, Offset)) {
+ } else if (PointerSize == 64
+ ? SelectADDRsi64(Op1.getNode(), Op1, Base, Offset)
+ : SelectADDRsi(Op1.getNode(), Op1, Base, Offset)) {
switch (N->getOpcode()) {
default:
return false;
@@ -1090,9 +1107,10 @@ bool NVPTXDAGToDAGISel::tryLoadVector(SD
getI32Imm(VecType, DL), getI32Imm(FromType, DL),
getI32Imm(FromTypeWidth, DL), Base, Offset, Chain };
LD = CurDAG->getMachineNode(Opcode.getValue(), DL, N->getVTList(), Ops);
- } else if (TM.is64Bit() ? SelectADDRri64(Op1.getNode(), Op1, Base, Offset)
- : SelectADDRri(Op1.getNode(), Op1, Base, Offset)) {
- if (TM.is64Bit()) {
+ } else if (PointerSize == 64
+ ? SelectADDRri64(Op1.getNode(), Op1, Base, Offset)
+ : SelectADDRri(Op1.getNode(), Op1, Base, Offset)) {
+ if (PointerSize == 64) {
switch (N->getOpcode()) {
default:
return false;
@@ -1140,7 +1158,7 @@ bool NVPTXDAGToDAGISel::tryLoadVector(SD
LD = CurDAG->getMachineNode(Opcode.getValue(), DL, N->getVTList(), Ops);
} else {
- if (TM.is64Bit()) {
+ if (PointerSize == 64) {
switch (N->getOpcode()) {
default:
return false;
@@ -1685,14 +1703,16 @@ bool NVPTXDAGToDAGISel::tryStore(SDNode
return false;
// Address Space Setting
- unsigned int codeAddrSpace = getCodeAddrSpace(ST);
+ unsigned int CodeAddrSpace = getCodeAddrSpace(ST);
+ unsigned int PointerSize =
+ CurDAG->getDataLayout().getPointerSizeInBits(ST->getAddressSpace());
// Volatile Setting
// - .volatile is only availalble for .global and .shared
bool isVolatile = ST->isVolatile();
- if (codeAddrSpace != NVPTX::PTXLdStInstCode::GLOBAL &&
- codeAddrSpace != NVPTX::PTXLdStInstCode::SHARED &&
- codeAddrSpace != NVPTX::PTXLdStInstCode::GENERIC)
+ if (CodeAddrSpace != NVPTX::PTXLdStInstCode::GLOBAL &&
+ CodeAddrSpace != NVPTX::PTXLdStInstCode::SHARED &&
+ CodeAddrSpace != NVPTX::PTXLdStInstCode::GENERIC)
isVolatile = false;
// Vector Setting
@@ -1735,12 +1755,12 @@ bool NVPTXDAGToDAGISel::tryStore(SDNode
if (!Opcode)
return false;
SDValue Ops[] = { N1, getI32Imm(isVolatile, dl),
- getI32Imm(codeAddrSpace, dl), getI32Imm(vecType, dl),
+ getI32Imm(CodeAddrSpace, dl), getI32Imm(vecType, dl),
getI32Imm(toType, dl), getI32Imm(toTypeWidth, dl), Addr,
Chain };
NVPTXST = CurDAG->getMachineNode(Opcode.getValue(), dl, MVT::Other, Ops);
- } else if (TM.is64Bit() ? SelectADDRsi64(N2.getNode(), N2, Base, Offset)
- : SelectADDRsi(N2.getNode(), N2, Base, Offset)) {
+ } else if (PointerSize == 64 ? SelectADDRsi64(N2.getNode(), N2, Base, Offset)
+ : SelectADDRsi(N2.getNode(), N2, Base, Offset)) {
Opcode = pickOpcodeForVT(SourceVT, NVPTX::ST_i8_asi, NVPTX::ST_i16_asi,
NVPTX::ST_i32_asi, NVPTX::ST_i64_asi,
NVPTX::ST_f16_asi, NVPTX::ST_f16x2_asi,
@@ -1748,13 +1768,13 @@ bool NVPTXDAGToDAGISel::tryStore(SDNode
if (!Opcode)
return false;
SDValue Ops[] = { N1, getI32Imm(isVolatile, dl),
- getI32Imm(codeAddrSpace, dl), getI32Imm(vecType, dl),
+ getI32Imm(CodeAddrSpace, dl), getI32Imm(vecType, dl),
getI32Imm(toType, dl), getI32Imm(toTypeWidth, dl), Base,
Offset, Chain };
NVPTXST = CurDAG->getMachineNode(Opcode.getValue(), dl, MVT::Other, Ops);
- } else if (TM.is64Bit() ? SelectADDRri64(N2.getNode(), N2, Base, Offset)
- : SelectADDRri(N2.getNode(), N2, Base, Offset)) {
- if (TM.is64Bit())
+ } else if (PointerSize == 64 ? SelectADDRri64(N2.getNode(), N2, Base, Offset)
+ : SelectADDRri(N2.getNode(), N2, Base, Offset)) {
+ if (PointerSize == 64)
Opcode = pickOpcodeForVT(
SourceVT, NVPTX::ST_i8_ari_64, NVPTX::ST_i16_ari_64,
NVPTX::ST_i32_ari_64, NVPTX::ST_i64_ari_64, NVPTX::ST_f16_ari_64,
@@ -1768,12 +1788,12 @@ bool NVPTXDAGToDAGISel::tryStore(SDNode
return false;
SDValue Ops[] = { N1, getI32Imm(isVolatile, dl),
- getI32Imm(codeAddrSpace, dl), getI32Imm(vecType, dl),
+ getI32Imm(CodeAddrSpace, dl), getI32Imm(vecType, dl),
getI32Imm(toType, dl), getI32Imm(toTypeWidth, dl), Base,
Offset, Chain };
NVPTXST = CurDAG->getMachineNode(Opcode.getValue(), dl, MVT::Other, Ops);
} else {
- if (TM.is64Bit())
+ if (PointerSize == 64)
Opcode =
pickOpcodeForVT(SourceVT, NVPTX::ST_i8_areg_64, NVPTX::ST_i16_areg_64,
NVPTX::ST_i32_areg_64, NVPTX::ST_i64_areg_64,
@@ -1787,7 +1807,7 @@ bool NVPTXDAGToDAGISel::tryStore(SDNode
if (!Opcode)
return false;
SDValue Ops[] = { N1, getI32Imm(isVolatile, dl),
- getI32Imm(codeAddrSpace, dl), getI32Imm(vecType, dl),
+ getI32Imm(CodeAddrSpace, dl), getI32Imm(vecType, dl),
getI32Imm(toType, dl), getI32Imm(toTypeWidth, dl), N2,
Chain };
NVPTXST = CurDAG->getMachineNode(Opcode.getValue(), dl, MVT::Other, Ops);
@@ -1816,11 +1836,12 @@ bool NVPTXDAGToDAGISel::tryStoreVector(S
// Address Space Setting
unsigned CodeAddrSpace = getCodeAddrSpace(MemSD);
-
if (CodeAddrSpace == NVPTX::PTXLdStInstCode::CONSTANT) {
report_fatal_error("Cannot store to pointer that points to constant "
"memory space");
}
+ unsigned int PointerSize =
+ CurDAG->getDataLayout().getPointerSizeInBits(MemSD->getAddressSpace());
// Volatile Setting
// - .volatile is only availalble for .global and .shared
@@ -1901,8 +1922,8 @@ bool NVPTXDAGToDAGISel::tryStoreVector(S
break;
}
StOps.push_back(Addr);
- } else if (TM.is64Bit() ? SelectADDRsi64(N2.getNode(), N2, Base, Offset)
- : SelectADDRsi(N2.getNode(), N2, Base, Offset)) {
+ } else if (PointerSize == 64 ? SelectADDRsi64(N2.getNode(), N2, Base, Offset)
+ : SelectADDRsi(N2.getNode(), N2, Base, Offset)) {
switch (N->getOpcode()) {
default:
return false;
@@ -1923,9 +1944,9 @@ bool NVPTXDAGToDAGISel::tryStoreVector(S
}
StOps.push_back(Base);
StOps.push_back(Offset);
- } else if (TM.is64Bit() ? SelectADDRri64(N2.getNode(), N2, Base, Offset)
- : SelectADDRri(N2.getNode(), N2, Base, Offset)) {
- if (TM.is64Bit()) {
+ } else if (PointerSize == 64 ? SelectADDRri64(N2.getNode(), N2, Base, Offset)
+ : SelectADDRri(N2.getNode(), N2, Base, Offset)) {
+ if (PointerSize == 64) {
switch (N->getOpcode()) {
default:
return false;
@@ -1968,7 +1989,7 @@ bool NVPTXDAGToDAGISel::tryStoreVector(S
StOps.push_back(Base);
StOps.push_back(Offset);
} else {
- if (TM.is64Bit()) {
+ if (PointerSize == 64) {
switch (N->getOpcode()) {
default:
return false;
Modified: llvm/trunk/lib/Target/NVPTX/NVPTXISelDAGToDAG.h
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXISelDAGToDAG.h?rev=331941&r1=331940&r2=331941&view=diff
==============================================================================
--- llvm/trunk/lib/Target/NVPTX/NVPTXISelDAGToDAG.h (original)
+++ llvm/trunk/lib/Target/NVPTX/NVPTXISelDAGToDAG.h Wed May 9 16:46:19 2018
@@ -35,6 +35,7 @@ class LLVM_LIBRARY_VISIBILITY NVPTXDAGTo
bool useF32FTZ() const;
bool allowFMA() const;
bool allowUnsafeFPMath() const;
+ bool useShortPointers() const;
public:
explicit NVPTXDAGToDAGISel(NVPTXTargetMachine &tm,
Modified: llvm/trunk/lib/Target/NVPTX/NVPTXISelLowering.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXISelLowering.cpp?rev=331941&r1=331940&r2=331941&view=diff
==============================================================================
--- llvm/trunk/lib/Target/NVPTX/NVPTXISelLowering.cpp (original)
+++ llvm/trunk/lib/Target/NVPTX/NVPTXISelLowering.cpp Wed May 9 16:46:19 2018
@@ -1233,9 +1233,9 @@ SDValue NVPTXTargetLowering::getSqrtEsti
SDValue
NVPTXTargetLowering::LowerGlobalAddress(SDValue Op, SelectionDAG &DAG) const {
SDLoc dl(Op);
- const GlobalValue *GV = cast<GlobalAddressSDNode>(Op)->getGlobal();
- auto PtrVT = getPointerTy(DAG.getDataLayout());
- Op = DAG.getTargetGlobalAddress(GV, dl, PtrVT);
+ const GlobalAddressSDNode *GAN = cast<GlobalAddressSDNode>(Op);
+ auto PtrVT = getPointerTy(DAG.getDataLayout(), GAN->getAddressSpace());
+ Op = DAG.getTargetGlobalAddress(GAN->getGlobal(), dl, PtrVT);
return DAG.getNode(NVPTXISD::Wrapper, dl, PtrVT, Op);
}
Modified: llvm/trunk/lib/Target/NVPTX/NVPTXInstrInfo.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXInstrInfo.td?rev=331941&r1=331940&r2=331941&view=diff
==============================================================================
--- llvm/trunk/lib/Target/NVPTX/NVPTXInstrInfo.td (original)
+++ llvm/trunk/lib/Target/NVPTX/NVPTXInstrInfo.td Wed May 9 16:46:19 2018
@@ -147,6 +147,7 @@ def hasPTX61 : Predicate<"Subtarget->get
def hasSM30 : Predicate<"Subtarget->getSmVersion() >= 30">;
def hasSM70 : Predicate<"Subtarget->getSmVersion() >= 70">;
+def useShortPtr : Predicate<"useShortPointers()">;
def useFP16Math: Predicate<"Subtarget->allowFP16Math()">;
//===----------------------------------------------------------------------===//
Modified: llvm/trunk/lib/Target/NVPTX/NVPTXIntrinsics.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXIntrinsics.td?rev=331941&r1=331940&r2=331941&view=diff
==============================================================================
--- llvm/trunk/lib/Target/NVPTX/NVPTXIntrinsics.td (original)
+++ llvm/trunk/lib/Target/NVPTX/NVPTXIntrinsics.td Wed May 9 16:46:19 2018
@@ -1937,6 +1937,12 @@ multiclass NG_TO_G<string Str, Intrinsic
def _yes_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
!strconcat("cvta.", Str, ".u64 \t$result, $src;"),
[(set Int64Regs:$result, (Intrin Int64Regs:$src))]>;
+ def _yes_6432 : NVPTXInst<(outs Int64Regs:$result), (ins Int32Regs:$src),
+ "{{ .reg .b64 %tmp;\n\t"
+ #" cvt.u64.u32 \t%tmp, $src;\n\t"
+ #" cvta." # Str # ".u64 \t$result, %tmp; }}",
+ [(set Int64Regs:$result, (Intrin Int32Regs:$src))]>,
+ Requires<[useShortPtr]>;
}
multiclass G_TO_NG<string Str, Intrinsic Intrin> {
@@ -1946,6 +1952,12 @@ multiclass G_TO_NG<string Str, Intrinsic
def _yes_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
!strconcat("cvta.to.", Str, ".u64 \t$result, $src;"),
[(set Int64Regs:$result, (Intrin Int64Regs:$src))]>;
+ def _yes_3264 : NVPTXInst<(outs Int32Regs:$result), (ins Int64Regs:$src),
+ "{{ .reg .b64 %tmp;\n\t"
+ #" cvta.to." # Str # ".u64 \t%tmp, $src;\n\t"
+ #" cvt.u32.u64 \t$result, %tmp; }}",
+ [(set Int32Regs:$result, (Intrin Int64Regs:$src))]>,
+ Requires<[useShortPtr]>;
}
defm cvta_local : NG_TO_G<"local", int_nvvm_ptr_local_to_gen>;
Modified: llvm/trunk/lib/Target/NVPTX/NVPTXSubtarget.h
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXSubtarget.h?rev=331941&r1=331940&r2=331941&view=diff
==============================================================================
--- llvm/trunk/lib/Target/NVPTX/NVPTXSubtarget.h (original)
+++ llvm/trunk/lib/Target/NVPTX/NVPTXSubtarget.h Wed May 9 16:46:19 2018
@@ -79,7 +79,6 @@ public:
bool hasImageHandles() const;
bool hasFP16Math() const { return SmVersion >= 53; }
bool allowFP16Math() const;
-
unsigned int getSmVersion() const { return SmVersion; }
std::string getTargetName() const { return TargetName; }
Modified: llvm/trunk/lib/Target/NVPTX/NVPTXTargetMachine.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXTargetMachine.cpp?rev=331941&r1=331940&r2=331941&view=diff
==============================================================================
--- llvm/trunk/lib/Target/NVPTX/NVPTXTargetMachine.cpp (original)
+++ llvm/trunk/lib/Target/NVPTX/NVPTXTargetMachine.cpp Wed May 9 16:46:19 2018
@@ -52,6 +52,12 @@ static cl::opt<bool> DisableRequireStruc
"unexpected regressions happen."),
cl::init(false), cl::Hidden);
+static cl::opt<bool> UseShortPointersOpt(
+ "nvptx-short-ptr",
+ cl::desc(
+ "Use 32-bit pointers for accessing const/local/shared address spaces."),
+ cl::init(false), cl::Hidden);
+
namespace llvm {
void initializeNVVMIntrRangePass(PassRegistry&);
@@ -83,11 +89,13 @@ extern "C" void LLVMInitializeNVPTXTarge
initializeNVPTXLowerAggrCopiesPass(PR);
}
-static std::string computeDataLayout(bool is64Bit) {
+static std::string computeDataLayout(bool is64Bit, bool UseShortPointers) {
std::string Ret = "e";
if (!is64Bit)
Ret += "-p:32:32";
+ else if (UseShortPointers)
+ Ret += "-p3:32:32-p4:32:32-p5:32:32";
Ret += "-i64:64-i128:128-v16:16-v32:32-n16:32:64";
@@ -108,9 +116,11 @@ NVPTXTargetMachine::NVPTXTargetMachine(c
CodeGenOpt::Level OL, bool is64bit)
// The pic relocation model is used regardless of what the client has
// specified, as it is the only relocation model currently supported.
- : LLVMTargetMachine(T, computeDataLayout(is64bit), TT, CPU, FS, Options,
- Reloc::PIC_, getEffectiveCodeModel(CM), OL),
- is64bit(is64bit), TLOF(llvm::make_unique<NVPTXTargetObjectFile>()),
+ : LLVMTargetMachine(T, computeDataLayout(is64bit, UseShortPointersOpt), TT,
+ CPU, FS, Options, Reloc::PIC_,
+ getEffectiveCodeModel(CM), OL),
+ is64bit(is64bit), UseShortPointers(UseShortPointersOpt),
+ TLOF(llvm::make_unique<NVPTXTargetObjectFile>()),
Subtarget(TT, CPU, FS, *this) {
if (TT.getOS() == Triple::NVCL)
drvInterface = NVPTX::NVCL;
Modified: llvm/trunk/lib/Target/NVPTX/NVPTXTargetMachine.h
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXTargetMachine.h?rev=331941&r1=331940&r2=331941&view=diff
==============================================================================
--- llvm/trunk/lib/Target/NVPTX/NVPTXTargetMachine.h (original)
+++ llvm/trunk/lib/Target/NVPTX/NVPTXTargetMachine.h Wed May 9 16:46:19 2018
@@ -26,6 +26,8 @@ namespace llvm {
///
class NVPTXTargetMachine : public LLVMTargetMachine {
bool is64bit;
+ // Use 32-bit pointers for accessing const/local/short AS.
+ bool UseShortPointers;
std::unique_ptr<TargetLoweringObjectFile> TLOF;
NVPTX::DrvInterface drvInterface;
NVPTXSubtarget Subtarget;
@@ -45,6 +47,7 @@ public:
}
const NVPTXSubtarget *getSubtargetImpl() const { return &Subtarget; }
bool is64Bit() const { return is64bit; }
+ bool useShortPointers() const { return UseShortPointers; }
NVPTX::DrvInterface getDrvInterface() const { return drvInterface; }
ManagedStringPool *getManagedStrPool() const {
return const_cast<ManagedStringPool *>(&ManagedStrPool);
Modified: llvm/trunk/test/CodeGen/NVPTX/addrspacecast.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/NVPTX/addrspacecast.ll?rev=331941&r1=331940&r2=331941&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/NVPTX/addrspacecast.ll (original)
+++ llvm/trunk/test/CodeGen/NVPTX/addrspacecast.ll Wed May 9 16:46:19 2018
@@ -1,97 +1,96 @@
-; RUN: llc -O0 < %s -march=nvptx -mcpu=sm_20 | FileCheck %s -check-prefix=PTX32
-; RUN: llc -O0 < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s -check-prefix=PTX64
+; RUN: llc -O0 < %s -march=nvptx -mcpu=sm_20 | FileCheck %s -check-prefixes=ALL,CLS32,G32
+; RUN: llc -O0 < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s -check-prefixes=ALL,NOPTRCONV,CLS64,G64
+; RUN: llc -O0 < %s -march=nvptx64 -mcpu=sm_20 --nvptx-short-ptr| FileCheck %s -check-prefixes=ALL,PTRCONV,CLS64,G64
+; ALL-LABEL: conv1
define i32 @conv1(i32 addrspace(1)* %ptr) {
-; PTX32: conv1
-; PTX32: cvta.global.u32
-; PTX32: ld.u32
-; PTX64: conv1
-; PTX64: cvta.global.u64
-; PTX64: ld.u32
+; G32: cvta.global.u32
+; ALL-NOT: cvt.u64.u32
+; G64: cvta.global.u64
+; ALL: ld.u32
%genptr = addrspacecast i32 addrspace(1)* %ptr to i32*
%val = load i32, i32* %genptr
ret i32 %val
}
+; ALL-LABEL: conv2
define i32 @conv2(i32 addrspace(3)* %ptr) {
-; PTX32: conv2
-; PTX32: cvta.shared.u32
-; PTX32: ld.u32
-; PTX64: conv2
-; PTX64: cvta.shared.u64
-; PTX64: ld.u32
+; CLS32: cvta.shared.u32
+; PTRCONV: cvt.u64.u32
+; NOPTRCONV-NOT: cvt.u64.u32
+; CLS64: cvta.shared.u64
+; ALL: ld.u32
%genptr = addrspacecast i32 addrspace(3)* %ptr to i32*
%val = load i32, i32* %genptr
ret i32 %val
}
+; ALL-LABEL: conv3
define i32 @conv3(i32 addrspace(4)* %ptr) {
-; PTX32: conv3
-; PTX32: cvta.const.u32
-; PTX32: ld.u32
-; PTX64: conv3
-; PTX64: cvta.const.u64
-; PTX64: ld.u32
+; CLS32: cvta.const.u32
+; PTRCONV: cvt.u64.u32
+; NOPTRCONV-NOT: cvt.u64.u32
+; CLS64: cvta.const.u64
+; ALL: ld.u32
%genptr = addrspacecast i32 addrspace(4)* %ptr to i32*
%val = load i32, i32* %genptr
ret i32 %val
}
+; ALL-LABEL: conv4
define i32 @conv4(i32 addrspace(5)* %ptr) {
-; PTX32: conv4
-; PTX32: cvta.local.u32
-; PTX32: ld.u32
-; PTX64: conv4
-; PTX64: cvta.local.u64
-; PTX64: ld.u32
+; CLS32: cvta.local.u32
+; PTRCONV: cvt.u64.u32
+; NOPTRCONV-NOT: cvt.u64.u32
+; CLS64: cvta.local.u64
+; ALL: ld.u32
%genptr = addrspacecast i32 addrspace(5)* %ptr to i32*
%val = load i32, i32* %genptr
ret i32 %val
}
+; ALL-LABEL: conv5
define i32 @conv5(i32* %ptr) {
-; PTX32: conv5
-; PTX32: cvta.to.global.u32
-; PTX32: ld.global.u32
-; PTX64: conv5
-; PTX64: cvta.to.global.u64
-; PTX64: ld.global.u32
+; CLS32: cvta.to.global.u32
+; ALL-NOT: cvt.u64.u32
+; CLS64: cvta.to.global.u64
+; ALL: ld.global.u32
%specptr = addrspacecast i32* %ptr to i32 addrspace(1)*
%val = load i32, i32 addrspace(1)* %specptr
ret i32 %val
}
+; ALL-LABEL: conv6
define i32 @conv6(i32* %ptr) {
-; PTX32: conv6
-; PTX32: cvta.to.shared.u32
-; PTX32: ld.shared.u32
-; PTX64: conv6
-; PTX64: cvta.to.shared.u64
-; PTX64: ld.shared.u32
+; CLS32: cvta.to.shared.u32
+; CLS64: cvta.to.shared.u64
+; PTRCONV: cvt.u32.u64
+; NOPTRCONV-NOT: cvt.u32.u64
+; ALL: ld.shared.u32
%specptr = addrspacecast i32* %ptr to i32 addrspace(3)*
%val = load i32, i32 addrspace(3)* %specptr
ret i32 %val
}
+; ALL-LABEL: conv7
define i32 @conv7(i32* %ptr) {
-; PTX32: conv7
-; PTX32: cvta.to.const.u32
-; PTX32: ld.const.u32
-; PTX64: conv7
-; PTX64: cvta.to.const.u64
-; PTX64: ld.const.u32
+; CLS32: cvta.to.const.u32
+; CLS64: cvta.to.const.u64
+; PTRCONV: cvt.u32.u64
+; NOPTRCONV-NOT: cvt.u32.u64
+; ALL: ld.const.u32
%specptr = addrspacecast i32* %ptr to i32 addrspace(4)*
%val = load i32, i32 addrspace(4)* %specptr
ret i32 %val
}
+; ALL-LABEL: conv8
define i32 @conv8(i32* %ptr) {
-; PTX32: conv8
-; PTX32: cvta.to.local.u32
-; PTX32: ld.local.u32
-; PTX64: conv8
-; PTX64: cvta.to.local.u64
-; PTX64: ld.local.u32
+; CLS32: cvta.to.local.u32
+; CLS64: cvta.to.local.u64
+; PTRCONV: cvt.u32.u64
+; NOPTRCONV-NOT: cvt.u32.u64
+; ALL: ld.local.u32
%specptr = addrspacecast i32* %ptr to i32 addrspace(5)*
%val = load i32, i32 addrspace(5)* %specptr
ret i32 %val
Modified: llvm/trunk/test/CodeGen/NVPTX/ld-addrspace.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/NVPTX/ld-addrspace.ll?rev=331941&r1=331940&r2=331941&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/NVPTX/ld-addrspace.ll (original)
+++ llvm/trunk/test/CodeGen/NVPTX/ld-addrspace.ll Wed May 9 16:46:19 2018
@@ -1,171 +1,160 @@
-; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefix=PTX32
-; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefix=PTX64
+; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefixes=ALL,G32,LS32
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefixes=ALL,G64,LS64
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 --nvptx-short-ptr | FileCheck %s --check-prefixes=G64,LS32
;; i8
define i8 @ld_global_i8(i8 addrspace(1)* %ptr) {
-; PTX32: ld.global.u8 %r{{[0-9]+}}, [%r{{[0-9]+}}]
-; PTX32: ret
-; PTX64: ld.global.u8 %r{{[0-9]+}}, [%rd{{[0-9]+}}]
-; PTX64: ret
+; ALL-LABEL: ld_global_i8
+; G32: ld.global.u8 %{{.*}}, [%r{{[0-9]+}}]
+; G64: ld.global.u8 %{{.*}}, [%rd{{[0-9]+}}]
+; ALL: ret
%a = load i8, i8 addrspace(1)* %ptr
ret i8 %a
}
-
define i8 @ld_shared_i8(i8 addrspace(3)* %ptr) {
-; PTX32: ld.shared.u8 %r{{[0-9]+}}, [%r{{[0-9]+}}]
-; PTX32: ret
-; PTX64: ld.shared.u8 %r{{[0-9]+}}, [%rd{{[0-9]+}}]
-; PTX64: ret
+; ALL-LABEL: ld_shared_i8
+; LS32: ld.shared.u8 %{{.*}}, [%r{{[0-9]+}}]
+; LS64: ld.shared.u8 %{{.*}}, [%rd{{[0-9]+}}]
+; ALL: ret
%a = load i8, i8 addrspace(3)* %ptr
ret i8 %a
}
-
define i8 @ld_local_i8(i8 addrspace(5)* %ptr) {
-; PTX32: ld.local.u8 %r{{[0-9]+}}, [%r{{[0-9]+}}]
-; PTX32: ret
-; PTX64: ld.local.u8 %r{{[0-9]+}}, [%rd{{[0-9]+}}]
-; PTX64: ret
+; ALL-LABEL: ld_local_i8
+; LS32: ld.local.u8 %{{.*}}, [%r{{[0-9]+}}]
+; LS64: ld.local.u8 %{{.*}}, [%rd{{[0-9]+}}]
+; ALL: ret
%a = load i8, i8 addrspace(5)* %ptr
ret i8 %a
}
;; i16
define i16 @ld_global_i16(i16 addrspace(1)* %ptr) {
-; PTX32: ld.global.u16 %r{{[0-9]+}}, [%r{{[0-9]+}}]
-; PTX32: ret
-; PTX64: ld.global.u16 %r{{[0-9]+}}, [%rd{{[0-9]+}}]
-; PTX64: ret
+; ALL-LABEL: ld_global_i16
+; G32: ld.global.u16 %{{.*}}, [%r{{[0-9]+}}]
+; G64: ld.global.u16 %{{.*}}, [%rd{{[0-9]+}}]
+; ALL: ret
%a = load i16, i16 addrspace(1)* %ptr
ret i16 %a
}
-
define i16 @ld_shared_i16(i16 addrspace(3)* %ptr) {
-; PTX32: ld.shared.u16 %r{{[0-9]+}}, [%r{{[0-9]+}}]
-; PTX32: ret
-; PTX64: ld.shared.u16 %r{{[0-9]+}}, [%rd{{[0-9]+}}]
-; PTX64: ret
+; ALL-LABEL: ld_shared_i16
+; LS32: ld.shared.u16 %{{.*}}, [%r{{[0-9]+}}]
+; LS64: ld.shared.u16 %{{.*}}, [%rd{{[0-9]+}}]
+; ALL: ret
%a = load i16, i16 addrspace(3)* %ptr
ret i16 %a
}
-
define i16 @ld_local_i16(i16 addrspace(5)* %ptr) {
-; PTX32: ld.local.u16 %r{{[0-9]+}}, [%r{{[0-9]+}}]
-; PTX32: ret
-; PTX64: ld.local.u16 %r{{[0-9]+}}, [%rd{{[0-9]+}}]
-; PTX64: ret
+; ALL-LABEL: ld_local_i16
+; LS32: ld.local.u16 %{{.*}}, [%r{{[0-9]+}}]
+; LS64: ld.local.u16 %{{.*}}, [%rd{{[0-9]+}}]
+; ALL: ret
%a = load i16, i16 addrspace(5)* %ptr
ret i16 %a
}
;; i32
define i32 @ld_global_i32(i32 addrspace(1)* %ptr) {
-; PTX32: ld.global.u32 %r{{[0-9]+}}, [%r{{[0-9]+}}]
-; PTX32: ret
-; PTX64: ld.global.u32 %r{{[0-9]+}}, [%rd{{[0-9]+}}]
-; PTX64: ret
+; ALL-LABEL: ld_global_i32
+; G32: ld.global.u32 %{{.*}}, [%r{{[0-9]+}}]
+; G64: ld.global.u32 %{{.*}}, [%rd{{[0-9]+}}]
+; ALL: ret
%a = load i32, i32 addrspace(1)* %ptr
ret i32 %a
}
-
define i32 @ld_shared_i32(i32 addrspace(3)* %ptr) {
-; PTX32: ld.shared.u32 %r{{[0-9]+}}, [%r{{[0-9]+}}]
-; PTX32: ret
-; PTX64: ld.shared.u32 %r{{[0-9]+}}, [%rd{{[0-9]+}}]
+; ALL-LABEL: ld_shared_i32
+; LS32: ld.shared.u32 %{{.*}}, [%r{{[0-9]+}}]
+; LS64: ld.shared.u32 %{{.*}}, [%rd{{[0-9]+}}]
; PTX64: ret
%a = load i32, i32 addrspace(3)* %ptr
ret i32 %a
}
-
define i32 @ld_local_i32(i32 addrspace(5)* %ptr) {
-; PTX32: ld.local.u32 %r{{[0-9]+}}, [%r{{[0-9]+}}]
-; PTX32: ret
-; PTX64: ld.local.u32 %r{{[0-9]+}}, [%rd{{[0-9]+}}]
-; PTX64: ret
+; ALL-LABEL: ld_local_i32
+; LS32: ld.local.u32 %{{.*}}, [%r{{[0-9]+}}]
+; LS64: ld.local.u32 %{{.*}}, [%rd{{[0-9]+}}]
+; ALL: ret
%a = load i32, i32 addrspace(5)* %ptr
ret i32 %a
}
;; i64
define i64 @ld_global_i64(i64 addrspace(1)* %ptr) {
-; PTX32: ld.global.u64 %rd{{[0-9]+}}, [%r{{[0-9]+}}]
-; PTX32: ret
-; PTX64: ld.global.u64 %rd{{[0-9]+}}, [%rd{{[0-9]+}}]
-; PTX64: ret
+; ALL-LABEL: ld_global_i64
+; G32: ld.global.u64 %{{.*}}, [%r{{[0-9]+}}]
+; G64: ld.global.u64 %{{.*}}, [%rd{{[0-9]+}}]
+; ALL: ret
%a = load i64, i64 addrspace(1)* %ptr
ret i64 %a
}
-
define i64 @ld_shared_i64(i64 addrspace(3)* %ptr) {
-; PTX32: ld.shared.u64 %rd{{[0-9]+}}, [%r{{[0-9]+}}]
-; PTX32: ret
-; PTX64: ld.shared.u64 %rd{{[0-9]+}}, [%rd{{[0-9]+}}]
-; PTX64: ret
+; ALL-LABEL: ld_shared_i64
+; LS32: ld.shared.u64 %{{.*}}, [%r{{[0-9]+}}]
+; LS64: ld.shared.u64 %{{.*}}, [%rd{{[0-9]+}}]
+; ALL: ret
%a = load i64, i64 addrspace(3)* %ptr
ret i64 %a
}
-
define i64 @ld_local_i64(i64 addrspace(5)* %ptr) {
-; PTX32: ld.local.u64 %rd{{[0-9]+}}, [%r{{[0-9]+}}]
-; PTX32: ret
-; PTX64: ld.local.u64 %rd{{[0-9]+}}, [%rd{{[0-9]+}}]
-; PTX64: ret
+; ALL-LABEL: ld_local_i64
+; LS32: ld.local.u64 %{{.*}}, [%r{{[0-9]+}}]
+; LS64: ld.local.u64 %{{.*}}, [%rd{{[0-9]+}}]
+; ALL: ret
%a = load i64, i64 addrspace(5)* %ptr
ret i64 %a
}
;; f32
define float @ld_global_f32(float addrspace(1)* %ptr) {
-; PTX32: ld.global.f32 %f{{[0-9]+}}, [%r{{[0-9]+}}]
-; PTX32: ret
-; PTX64: ld.global.f32 %f{{[0-9]+}}, [%rd{{[0-9]+}}]
-; PTX64: ret
+; ALL-LABEL: ld_global_f32
+; G32: ld.global.f32 %{{.*}}, [%r{{[0-9]+}}]
+; G64: ld.global.f32 %{{.*}}, [%rd{{[0-9]+}}]
+; ALL: ret
%a = load float, float addrspace(1)* %ptr
ret float %a
}
-
define float @ld_shared_f32(float addrspace(3)* %ptr) {
-; PTX32: ld.shared.f32 %f{{[0-9]+}}, [%r{{[0-9]+}}]
-; PTX32: ret
-; PTX64: ld.shared.f32 %f{{[0-9]+}}, [%rd{{[0-9]+}}]
-; PTX64: ret
+; ALL-LABEL: ld_shared_f32
+; LS32: ld.shared.f32 %{{.*}}, [%r{{[0-9]+}}]
+; LS64: ld.shared.f32 %{{.*}}, [%rd{{[0-9]+}}]
+; ALL: ret
%a = load float, float addrspace(3)* %ptr
ret float %a
}
-
define float @ld_local_f32(float addrspace(5)* %ptr) {
-; PTX32: ld.local.f32 %f{{[0-9]+}}, [%r{{[0-9]+}}]
-; PTX32: ret
-; PTX64: ld.local.f32 %f{{[0-9]+}}, [%rd{{[0-9]+}}]
-; PTX64: ret
+; ALL-LABEL: ld_local_f32
+; LS32: ld.local.f32 %{{.*}}, [%r{{[0-9]+}}]
+; LS64: ld.local.f32 %{{.*}}, [%rd{{[0-9]+}}]
+; ALL: ret
%a = load float, float addrspace(5)* %ptr
ret float %a
}
;; f64
define double @ld_global_f64(double addrspace(1)* %ptr) {
-; PTX32: ld.global.f64 %fd{{[0-9]+}}, [%r{{[0-9]+}}]
-; PTX32: ret
-; PTX64: ld.global.f64 %fd{{[0-9]+}}, [%rd{{[0-9]+}}]
-; PTX64: ret
+; ALL-LABEL: ld_global_f64
+; G32: ld.global.f64 %{{.*}}, [%r{{[0-9]+}}]
+; G64: ld.global.f64 %{{.*}}, [%rd{{[0-9]+}}]
+; ALL: ret
%a = load double, double addrspace(1)* %ptr
ret double %a
}
-
define double @ld_shared_f64(double addrspace(3)* %ptr) {
-; PTX32: ld.shared.f64 %fd{{[0-9]+}}, [%r{{[0-9]+}}]
-; PTX32: ret
-; PTX64: ld.shared.f64 %fd{{[0-9]+}}, [%rd{{[0-9]+}}]
-; PTX64: ret
+; ALL-LABEL: ld_shared_f64
+; LS32: ld.shared.f64 %{{.*}}, [%r{{[0-9]+}}]
+; LS64: ld.shared.f64 %{{.*}}, [%rd{{[0-9]+}}]
+; ALL: ret
%a = load double, double addrspace(3)* %ptr
ret double %a
}
-
define double @ld_local_f64(double addrspace(5)* %ptr) {
-; PTX32: ld.local.f64 %fd{{[0-9]+}}, [%r{{[0-9]+}}]
-; PTX32: ret
-; PTX64: ld.local.f64 %fd{{[0-9]+}}, [%rd{{[0-9]+}}]
-; PTX64: ret
+; ALL-LABEL: ld_local_f64
+; LS32: ld.local.f64 %{{.*}}, [%r{{[0-9]+}}]
+; LS64: ld.local.f64 %{{.*}}, [%rd{{[0-9]+}}]
+; ALL: ret
%a = load double, double addrspace(5)* %ptr
ret double %a
}
Modified: llvm/trunk/test/CodeGen/NVPTX/st-addrspace.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/NVPTX/st-addrspace.ll?rev=331941&r1=331940&r2=331941&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/NVPTX/st-addrspace.ll (original)
+++ llvm/trunk/test/CodeGen/NVPTX/st-addrspace.ll Wed May 9 16:46:19 2018
@@ -1,177 +1,160 @@
-; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefix=PTX32
-; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefix=PTX64
+; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefixes=ALL,G32,LS32
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefixes=ALL,G64,LS64
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 --nvptx-short-ptr | FileCheck %s --check-prefixes=G64,LS32
;; i8
-
+; ALL-LABEL: st_global_i8
define void @st_global_i8(i8 addrspace(1)* %ptr, i8 %a) {
-; PTX32: st.global.u8 [%r{{[0-9]+}}], %rs{{[0-9]+}}
-; PTX32: ret
-; PTX64: st.global.u8 [%rd{{[0-9]+}}], %rs{{[0-9]+}}
-; PTX64: ret
+; G32: st.global.u8 [%r{{[0-9]+}}], %rs{{[0-9]+}}
+; G64: st.global.u8 [%rd{{[0-9]+}}], %rs{{[0-9]+}}
+; ALL: ret
store i8 %a, i8 addrspace(1)* %ptr
ret void
}
-
+; ALL-LABEL: st_shared_i8
define void @st_shared_i8(i8 addrspace(3)* %ptr, i8 %a) {
-; PTX32: st.shared.u8 [%r{{[0-9]+}}], %rs{{[0-9]+}}
-; PTX32: ret
-; PTX64: st.shared.u8 [%rd{{[0-9]+}}], %rs{{[0-9]+}}
-; PTX64: ret
+; LS32: st.shared.u8 [%r{{[0-9]+}}], %rs{{[0-9]+}}
+; LS64: st.shared.u8 [%rd{{[0-9]+}}], %rs{{[0-9]+}}
+; ALL: ret
store i8 %a, i8 addrspace(3)* %ptr
ret void
}
-
+; ALL-LABEL: st_local_i8
define void @st_local_i8(i8 addrspace(5)* %ptr, i8 %a) {
-; PTX32: st.local.u8 [%r{{[0-9]+}}], %rs{{[0-9]+}}
-; PTX32: ret
-; PTX64: st.local.u8 [%rd{{[0-9]+}}], %rs{{[0-9]+}}
-; PTX64: ret
+; LS32: st.local.u8 [%r{{[0-9]+}}], %rs{{[0-9]+}}
+; LS64: st.local.u8 [%rd{{[0-9]+}}], %rs{{[0-9]+}}
+; ALL: ret
store i8 %a, i8 addrspace(5)* %ptr
ret void
}
;; i16
-
+; ALL-LABEL: st_global_i16
define void @st_global_i16(i16 addrspace(1)* %ptr, i16 %a) {
-; PTX32: st.global.u16 [%r{{[0-9]+}}], %rs{{[0-9]+}}
-; PTX32: ret
-; PTX64: st.global.u16 [%rd{{[0-9]+}}], %rs{{[0-9]+}}
-; PTX64: ret
+; G32: st.global.u16 [%r{{[0-9]+}}], %rs{{[0-9]+}}
+; G64: st.global.u16 [%rd{{[0-9]+}}], %rs{{[0-9]+}}
+; ALL: ret
store i16 %a, i16 addrspace(1)* %ptr
ret void
}
-
+; ALL-LABEL: st_shared_i16
define void @st_shared_i16(i16 addrspace(3)* %ptr, i16 %a) {
-; PTX32: st.shared.u16 [%r{{[0-9]+}}], %rs{{[0-9]+}}
-; PTX32: ret
-; PTX64: st.shared.u16 [%rd{{[0-9]+}}], %rs{{[0-9]+}}
-; PTX64: ret
+; LS32: st.shared.u16 [%r{{[0-9]+}}], %rs{{[0-9]+}}
+; LS64: st.shared.u16 [%rd{{[0-9]+}}], %rs{{[0-9]+}}
+; ALL: ret
store i16 %a, i16 addrspace(3)* %ptr
ret void
}
-
+; ALL-LABEL: st_local_i16
define void @st_local_i16(i16 addrspace(5)* %ptr, i16 %a) {
-; PTX32: st.local.u16 [%r{{[0-9]+}}], %rs{{[0-9]+}}
-; PTX32: ret
-; PTX64: st.local.u16 [%rd{{[0-9]+}}], %rs{{[0-9]+}}
-; PTX64: ret
+; LS32: st.local.u16 [%r{{[0-9]+}}], %rs{{[0-9]+}}
+; LS64: st.local.u16 [%rd{{[0-9]+}}], %rs{{[0-9]+}}
+; ALL: ret
store i16 %a, i16 addrspace(5)* %ptr
ret void
}
;; i32
-
+; ALL-LABEL: st_global_i32
define void @st_global_i32(i32 addrspace(1)* %ptr, i32 %a) {
-; PTX32: st.global.u32 [%r{{[0-9]+}}], %r{{[0-9]+}}
-; PTX32: ret
-; PTX64: st.global.u32 [%rd{{[0-9]+}}], %r{{[0-9]+}}
-; PTX64: ret
+; G32: st.global.u32 [%r{{[0-9]+}}], %r{{[0-9]+}}
+; G64: st.global.u32 [%rd{{[0-9]+}}], %r{{[0-9]+}}
+; ALL: ret
store i32 %a, i32 addrspace(1)* %ptr
ret void
}
-
+; ALL-LABEL: st_shared_i32
define void @st_shared_i32(i32 addrspace(3)* %ptr, i32 %a) {
-; PTX32: st.shared.u32 [%r{{[0-9]+}}], %r{{[0-9]+}}
-; PTX32: ret
-; PTX64: st.shared.u32 [%rd{{[0-9]+}}], %r{{[0-9]+}}
+; LS32: st.shared.u32 [%r{{[0-9]+}}], %r{{[0-9]+}}
+; LS64: st.shared.u32 [%rd{{[0-9]+}}], %r{{[0-9]+}}
; PTX64: ret
store i32 %a, i32 addrspace(3)* %ptr
ret void
}
-
+; ALL-LABEL: st_local_i32
define void @st_local_i32(i32 addrspace(5)* %ptr, i32 %a) {
-; PTX32: st.local.u32 [%r{{[0-9]+}}], %r{{[0-9]+}}
-; PTX32: ret
-; PTX64: st.local.u32 [%rd{{[0-9]+}}], %r{{[0-9]+}}
-; PTX64: ret
+; LS32: st.local.u32 [%r{{[0-9]+}}], %r{{[0-9]+}}
+; LS64: st.local.u32 [%rd{{[0-9]+}}], %r{{[0-9]+}}
+; ALL: ret
store i32 %a, i32 addrspace(5)* %ptr
ret void
}
;; i64
-
+; ALL-LABEL: st_global_i64
define void @st_global_i64(i64 addrspace(1)* %ptr, i64 %a) {
-; PTX32: st.global.u64 [%r{{[0-9]+}}], %rd{{[0-9]+}}
-; PTX32: ret
-; PTX64: st.global.u64 [%rd{{[0-9]+}}], %rd{{[0-9]+}}
-; PTX64: ret
+; G32: st.global.u64 [%r{{[0-9]+}}], %rd{{[0-9]+}}
+; G64: st.global.u64 [%rd{{[0-9]+}}], %rd{{[0-9]+}}
+; ALL: ret
store i64 %a, i64 addrspace(1)* %ptr
ret void
}
-
+; ALL-LABEL: st_shared_i64
define void @st_shared_i64(i64 addrspace(3)* %ptr, i64 %a) {
-; PTX32: st.shared.u64 [%r{{[0-9]+}}], %rd{{[0-9]+}}
-; PTX32: ret
-; PTX64: st.shared.u64 [%rd{{[0-9]+}}], %rd{{[0-9]+}}
-; PTX64: ret
+; LS32: st.shared.u64 [%r{{[0-9]+}}], %rd{{[0-9]+}}
+; LS64: st.shared.u64 [%rd{{[0-9]+}}], %rd{{[0-9]+}}
+; ALL: ret
store i64 %a, i64 addrspace(3)* %ptr
ret void
}
-
+; ALL-LABEL: st_local_i64
define void @st_local_i64(i64 addrspace(5)* %ptr, i64 %a) {
-; PTX32: st.local.u64 [%r{{[0-9]+}}], %rd{{[0-9]+}}
-; PTX32: ret
-; PTX64: st.local.u64 [%rd{{[0-9]+}}], %rd{{[0-9]+}}
-; PTX64: ret
+; LS32: st.local.u64 [%r{{[0-9]+}}], %rd{{[0-9]+}}
+; LS64: st.local.u64 [%rd{{[0-9]+}}], %rd{{[0-9]+}}
+; ALL: ret
store i64 %a, i64 addrspace(5)* %ptr
ret void
}
;; f32
-
+; ALL-LABEL: st_global_f32
define void @st_global_f32(float addrspace(1)* %ptr, float %a) {
-; PTX32: st.global.f32 [%r{{[0-9]+}}], %f{{[0-9]+}}
-; PTX32: ret
-; PTX64: st.global.f32 [%rd{{[0-9]+}}], %f{{[0-9]+}}
-; PTX64: ret
+; G32: st.global.f32 [%r{{[0-9]+}}], %f{{[0-9]+}}
+; G64: st.global.f32 [%rd{{[0-9]+}}], %f{{[0-9]+}}
+; ALL: ret
store float %a, float addrspace(1)* %ptr
ret void
}
-
+; ALL-LABEL: st_shared_f32
define void @st_shared_f32(float addrspace(3)* %ptr, float %a) {
-; PTX32: st.shared.f32 [%r{{[0-9]+}}], %f{{[0-9]+}}
-; PTX32: ret
-; PTX64: st.shared.f32 [%rd{{[0-9]+}}], %f{{[0-9]+}}
-; PTX64: ret
+; LS32: st.shared.f32 [%r{{[0-9]+}}], %f{{[0-9]+}}
+; LS64: st.shared.f32 [%rd{{[0-9]+}}], %f{{[0-9]+}}
+; ALL: ret
store float %a, float addrspace(3)* %ptr
ret void
}
-
+; ALL-LABEL: st_local_f32
define void @st_local_f32(float addrspace(5)* %ptr, float %a) {
-; PTX32: st.local.f32 [%r{{[0-9]+}}], %f{{[0-9]+}}
-; PTX32: ret
-; PTX64: st.local.f32 [%rd{{[0-9]+}}], %f{{[0-9]+}}
-; PTX64: ret
+; LS32: st.local.f32 [%r{{[0-9]+}}], %f{{[0-9]+}}
+; LS64: st.local.f32 [%rd{{[0-9]+}}], %f{{[0-9]+}}
+; ALL: ret
store float %a, float addrspace(5)* %ptr
ret void
}
;; f64
-
+; ALL-LABEL: st_global_f64
define void @st_global_f64(double addrspace(1)* %ptr, double %a) {
-; PTX32: st.global.f64 [%r{{[0-9]+}}], %fd{{[0-9]+}}
-; PTX32: ret
-; PTX64: st.global.f64 [%rd{{[0-9]+}}], %fd{{[0-9]+}}
-; PTX64: ret
+; G32: st.global.f64 [%r{{[0-9]+}}], %fd{{[0-9]+}}
+; G64: st.global.f64 [%rd{{[0-9]+}}], %fd{{[0-9]+}}
+; ALL: ret
store double %a, double addrspace(1)* %ptr
ret void
}
-
+; ALL-LABEL: st_shared_f64
define void @st_shared_f64(double addrspace(3)* %ptr, double %a) {
-; PTX32: st.shared.f64 [%r{{[0-9]+}}], %fd{{[0-9]+}}
-; PTX32: ret
-; PTX64: st.shared.f64 [%rd{{[0-9]+}}], %fd{{[0-9]+}}
-; PTX64: ret
+; LS32: st.shared.f64 [%r{{[0-9]+}}], %fd{{[0-9]+}}
+; LS64: st.shared.f64 [%rd{{[0-9]+}}], %fd{{[0-9]+}}
+; ALL: ret
store double %a, double addrspace(3)* %ptr
ret void
}
-
+; ALL-LABEL: st_local_f64
define void @st_local_f64(double addrspace(5)* %ptr, double %a) {
-; PTX32: st.local.f64 [%r{{[0-9]+}}], %fd{{[0-9]+}}
-; PTX32: ret
-; PTX64: st.local.f64 [%rd{{[0-9]+}}], %fd{{[0-9]+}}
-; PTX64: ret
+; LS32: st.local.f64 [%r{{[0-9]+}}], %fd{{[0-9]+}}
+; LS64: st.local.f64 [%rd{{[0-9]+}}], %fd{{[0-9]+}}
+; ALL: ret
store double %a, double addrspace(5)* %ptr
ret void
}
More information about the llvm-commits
mailing list