[NVPTX] Added a feature to use short pointers for const/local/shared AS.

[tra]

Summary

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.

[jlebar]

Ben, are you comfortable reviewing this?

llvm/lib/Target/NVPTX/NVPTX.td

78 ↗

(On Diff #144215)

Should we s/short/32-bit/, or do you want to leave open the possibility of using shorter pointers?

llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp

738 ↗

(On Diff #144215)

s/ShortPointer /ShortPointers /?

Same below.

[tra]

UseShortPointer -> UseShortPointers

[tra]

Updated feature description.

[tra]

llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp

738 ↗

(On Diff #144215)

In this case we're dealing with just one pointer, while the subtarget can use any number of them. :-)
Oh, well, plural it is.

[bkramer]

llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp

855 ↗

(On Diff #144230)

There are 4 identical instances of the pointer size computation in this patch, it should really be pulled into its own function.

You can also consider getting the pointer size from the DataLayout instead. Might be cleaner.

llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp

1255 ↗

(On Diff #144230)

Why is this needed? The pointer returned by getPointerTy should be right wrt data layout.

[tra]

Use DataLayout to obtain pointer size.

[tra]

llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp

855 ↗

(On Diff #144230)

That's indeed the right thing to do here. I've updated the patch to use DataLayout for determining pointer size.

[bkramer]

lgtm

[arsenm]

NVPTX wasn't already doing this? As this is a module property, a subtarget feature can't be used for this. For example what happens for a global value?

[tra]

In D46147#1081632, @arsenm wrote:

NVPTX wasn't already doing this? As this is a module property, a subtarget feature can't be used for this. For example what happens for a global value?

No, NVPTX is always using 64-bit pointers for all address spaces and that's one of the common situations where it gives NVCC an edge.

It's not clear what concerns you about global values. Could you elaborate or give me an example?

[arsenm]

In D46147#1081660, @tra wrote:

In D46147#1081632, @arsenm wrote:

NVPTX wasn't already doing this? As this is a module property, a subtarget feature can't be used for this. For example what happens for a global value?

No, NVPTX is always using 64-bit pointers for all address spaces and that's one of the common situations where it gives NVCC an edge.

It's not clear what concerns you about global values. Could you elaborate or give me an example?

The subtarget is a function level property. The pointer size / datalayout is a module property and can't be different between functions. If you have multiple functions with different pointer sizes in a module (or no functions at all), then what pointer size ends up used for a global? The global itself has a size, and any ConstantExpr pointers in its initializer will also need to have a size.

[tra]

In D46147#1081662, @arsenm wrote:

The subtarget is a function level property. The pointer size / datalayout is a module property and can't be different between functions. If you have multiple functions with different pointer sizes in a module (or no functions at all), then what pointer size ends up used for a global? The global itself has a size, and any ConstantExpr pointers in its initializer will also need to have a size.

I think I see. So, someone could add "target-attributes" to a function and that may potentially override the features. It does not harm us in this case as we only set datalayout once when it's called during creation NVPTXTargetMachine so only the features supplied on command line will be affected. That said, I do see your point that subtarget is not the right place for this option. After a bit of digging I've found a curious comment in LLVMTargetMachine, so this approach is both used in other places and considered problemetic.

// FIXME: Having an MCSubtargetInfo on the target machine is a hack due
// to some backends having subtarget feature dependent module level
// code generation. This is similar to the hack in the AsmPrinter for
// module level assembly etc.
STI = TheTarget.createMCSubtargetInfo(getTargetTriple().str(), getTargetCPU(),
                                      getTargetFeatureString());

One way out of this is to add a separate command-line option to control this. Is there a better way to do it?

[tra]

Moved control of the feature from subtarget feature to --nvptx-short-ptr command-line option.

[tra]

@arsenm Will this do?

[bkramer]

Would it be possible to get rid of the flag entirely and make this controlled only by the data layout? That's defined on the module level.

[arsenm]

In D46147#1081951, @bkramer wrote:

Would it be possible to get rid of the flag entirely and make this controlled only by the data layout? That's defined on the module level.

Yes, ultimately the switch is just changing the datalayout. If this is just a short term thing for migration I think a global flag is fine. If you want to be able to swap these properly long term, this is a different arch name / TargetMachine.

[tra]

In D46147#1081990, @arsenm wrote:

In D46147#1081951, @bkramer wrote:

Would it be possible to get rid of the flag entirely and make this controlled only by the data layout? That's defined on the module level.

Yes, ultimately the switch is just changing the datalayout. If this is just a short term thing for migration I think a global flag is fine. If you want to be able to swap these properly long term, this is a different arch name / TargetMachine.

I'm still not quite sure how clang's setting of data layout in the module plays with LLVM's. When they are out of sync, LLVM throws an error which suggests that we currently can't just have arbitrary datalayout specified by the module, it must also match LLVM's idea of the data layout used by particular back-end. At least, that's my understanding how things currently work. I might be wrong.

In any case, I do want the short pointers to eventually become the default. However, it's a somewhat risky change, so I need an escape hatch to be able to test the changes without breaking everyone. I know of at least one issue in the split-GEP pass (will send a patch soon). There are likely more.

[arsenm]

In D46147#1087005, @tra wrote:

In D46147#1081990, @arsenm wrote:

In D46147#1081951, @bkramer wrote:

Would it be possible to get rid of the flag entirely and make this controlled only by the data layout? That's defined on the module level.

Yes, ultimately the switch is just changing the datalayout. If this is just a short term thing for migration I think a global flag is fine. If you want to be able to swap these properly long term, this is a different arch name / TargetMachine.

I'm still not quite sure how clang's setting of data layout in the module plays with LLVM's. When they are out of sync, LLVM throws an error which suggests that we currently can't just have arbitrary datalayout specified by the module, it must also match LLVM's idea of the data layout used by particular back-end. At least, that's my understanding how things currently work. I might be wrong.

In any case, I do want the short pointers to eventually become the default. However, it's a somewhat risky change, so I need an escape hatch to be able to test the changes without breaking everyone. I know of at least one issue in the split-GEP pass (will send a patch soon). There are likely more.

You need to repeat the datalayout in clang as the target expects

llvm/trunk/lib/Target/NVPTX/NVPTXISelDAGToDAG.h

class LLVM_LIBRARY_VISIBILITY NVPTXDAGToDAGISel : public SelectionDAGISel {
   // If true, generate mul.wide from sext and mul
   bool doMulWide;
 
   int getDivF32Level() const;
   bool usePrecSqrtF32() const;
   bool useF32FTZ() const;
   bool allowFMA() const;
   bool allowUnsafeFPMath() const;
+  bool useShortPointers() const;
 
 public:
   explicit NVPTXDAGToDAGISel(NVPTXTargetMachine &tm,
                              CodeGenOpt::Level   OptLevel);
 
   // Pass Name
   StringRef getPassName() const override {
     return "NVPTX DAG->DAG Pattern Instruction Selection";

llvm/trunk/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp

bool NVPTXDAGToDAGISel::allowFMA() const {
   return TL->allowFMA(*MF, OptLevel);
 }
 
 bool NVPTXDAGToDAGISel::allowUnsafeFPMath() const {
   const NVPTXTargetLowering *TL = Subtarget->getTargetLowering();
   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) {
 
   if (N->isMachineOpcode()) {
     N->setNodeId(-1);
     return; // Already selected.
   }
ReplaceNode(N, CurDAG->getMachineNode(NVPTX::texsurf_handles, SDLoc(N),
                                         MVT::i64, GlobalVal));
 }
 
 void NVPTXDAGToDAGISel::SelectAddrSpaceCast(SDNode *N) {
   SDValue Src = N->getOperand(0);
   AddrSpaceCastSDNode *CastN = cast<AddrSpaceCastSDNode>(N);
   unsigned SrcAddrSpace = CastN->getSrcAddressSpace();
   unsigned DstAddrSpace = CastN->getDestAddressSpace();
-
   assert(SrcAddrSpace != DstAddrSpace &&
          "addrspacecast must be between different address spaces");
 
   if (DstAddrSpace == ADDRESS_SPACE_GENERIC) {
     // Specific to generic
     unsigned Opc;
     switch (SrcAddrSpace) {
     default: report_fatal_error("Bad address space in addrspacecast");
     case ADDRESS_SPACE_GLOBAL:
       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),
                                           Src));
     return;
   } else {
     // Generic to specific
     if (SrcAddrSpace != 0)
       report_fatal_error("Cannot cast between two non-generic address spaces");
     unsigned Opc;
     switch (DstAddrSpace) {
     default: report_fatal_error("Bad address space in addrspacecast");
     case ADDRESS_SPACE_GLOBAL:
       Opc = TM.is64Bit() ? NVPTX::cvta_to_global_yes_64
                          : 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
                          : NVPTX::nvvm_ptr_gen_to_param;
       break;
     }
     ReplaceNode(N, CurDAG->getMachineNode(Opc, SDLoc(N), N->getValueType(0),
                                           Src));
bool NVPTXDAGToDAGISel::tryLoad(SDNode *N) {
   // do not support pre/post inc/dec
   if (LD->isIndexed())
     return false;
 
   if (!LoadedVT.isSimple())
     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
   //
   // Sign   : ISD::SEXTLOAD
   // Unsign : ISD::ZEXTLOAD, ISD::NON_EXTLOAD or ISD::EXTLOAD and the
   //          type is integer
   // Float  : ISD::NON_EXTLOAD or ISD::EXTLOAD and the type is float
bool NVPTXDAGToDAGISel::tryLoad(SDNode *N) {
 
   if (SelectDirectAddr(N1, Addr)) {
     Opcode = pickOpcodeForVT(
         TargetVT, NVPTX::LD_i8_avar, NVPTX::LD_i16_avar, NVPTX::LD_i32_avar,
         NVPTX::LD_i64_avar, NVPTX::LD_f16_avar, NVPTX::LD_f16x2_avar,
         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)
+  } 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)
+  } else if (PointerSize == 64 ? SelectADDRri64(N1.getNode(), N1, Base, Offset)
                                : SelectADDRri(N1.getNode(), N1, Base, Offset)) {
-    if (TM.is64Bit())
+    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,
           NVPTX::LD_f16x2_ari_64, NVPTX::LD_f32_ari_64, NVPTX::LD_f64_ari_64);
     else
       Opcode = pickOpcodeForVT(
           TargetVT, NVPTX::LD_i8_ari, NVPTX::LD_i16_ari, NVPTX::LD_i32_ari,
           NVPTX::LD_i64_ari, NVPTX::LD_f16_ari, NVPTX::LD_f16x2_ari,
           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,
           NVPTX::LD_f16x2_areg_64, NVPTX::LD_f32_areg_64,
           NVPTX::LD_f64_areg_64);
     else
       Opcode = pickOpcodeForVT(
           TargetVT, NVPTX::LD_i8_areg, NVPTX::LD_i16_areg, NVPTX::LD_i32_areg,
           NVPTX::LD_i64_areg, NVPTX::LD_f16_areg, NVPTX::LD_f16x2_areg,
           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,
                                      MVT::Other, Ops);
   }
 
   if (!NVPTXLD)
     return false;
bool NVPTXDAGToDAGISel::tryLoadVector(SDNode *N) {
   MemSDNode *MemSD = cast<MemSDNode>(N);
   EVT LoadedVT = MemSD->getMemoryVT();
 
   if (!LoadedVT.isSimple())
     return false;
 
   // 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();
   if (CodeAddrSpace != NVPTX::PTXLdStInstCode::GLOBAL &&
       CodeAddrSpace != NVPTX::PTXLdStInstCode::SHARED &&
       CodeAddrSpace != NVPTX::PTXLdStInstCode::GENERIC)
     IsVolatile = false;
 
case NVPTXISD::LoadV4:
       break;
     }
     if (!Opcode)
       return false;
     SDValue Ops[] = { getI32Imm(IsVolatile, DL), getI32Imm(CodeAddrSpace, DL),
                       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)
     switch (N->getOpcode()) {
     default:
       return false;
     case NVPTXISD::LoadV2:
       Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
                                NVPTX::LDV_i8_v2_asi, NVPTX::LDV_i16_v2_asi,
                                NVPTX::LDV_i32_v2_asi, NVPTX::LDV_i64_v2_asi,
                                NVPTX::LDV_f16_v2_asi, NVPTX::LDV_f16x2_v2_asi,
                                NVPTX::LDV_f32_v2_asi, NVPTX::LDV_f64_v2_asi);
       break;
     case NVPTXISD::LoadV4:
       Opcode =
           pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy, NVPTX::LDV_i8_v4_asi,
                           NVPTX::LDV_i16_v4_asi, NVPTX::LDV_i32_v4_asi, None,
                           NVPTX::LDV_f16_v4_asi, NVPTX::LDV_f16x2_v4_asi,
                           NVPTX::LDV_f32_v4_asi, None);
       break;
     }
     if (!Opcode)
       return false;
     SDValue Ops[] = { getI32Imm(IsVolatile, DL), getI32Imm(CodeAddrSpace, DL),
                       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)
+    if (PointerSize == 64) {
       switch (N->getOpcode()) {
       default:
         return false;
       case NVPTXISD::LoadV2:
         Opcode = pickOpcodeForVT(
             EltVT.getSimpleVT().SimpleTy, NVPTX::LDV_i8_v2_ari_64,
             NVPTX::LDV_i16_v2_ari_64, NVPTX::LDV_i32_v2_ari_64,
             NVPTX::LDV_i64_v2_ari_64, NVPTX::LDV_f16_v2_ari_64,
if (SelectDirectAddr(Op1, Addr)) {
     if (!Opcode)
       return false;
     SDValue Ops[] = { getI32Imm(IsVolatile, DL), getI32Imm(CodeAddrSpace, DL),
                       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()) {
+    if (PointerSize == 64) {
       switch (N->getOpcode()) {
       default:
         return false;
       case NVPTXISD::LoadV2:
         Opcode = pickOpcodeForVT(
             EltVT.getSimpleVT().SimpleTy, NVPTX::LDV_i8_v2_areg_64,
             NVPTX::LDV_i16_v2_areg_64, NVPTX::LDV_i32_v2_areg_64,
             NVPTX::LDV_i64_v2_areg_64, NVPTX::LDV_f16_v2_areg_64,
bool NVPTXDAGToDAGISel::tryStore(SDNode *N) {
   // do not support pre/post inc/dec
   if (ST->isIndexed())
     return false;
 
   if (!StoreVT.isSimple())
     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
   MVT SimpleVT = StoreVT.getSimpleVT();
   unsigned vecType = NVPTX::PTXLdStInstCode::Scalar;
 
   // Type Setting: toType + toTypeWidth
   // - for integer type, always use 'u'
bool NVPTXDAGToDAGISel::tryStore(SDNode *N) {
   if (SelectDirectAddr(N2, Addr)) {
     Opcode = pickOpcodeForVT(SourceVT, NVPTX::ST_i8_avar, NVPTX::ST_i16_avar,
                              NVPTX::ST_i32_avar, NVPTX::ST_i64_avar,
                              NVPTX::ST_f16_avar, NVPTX::ST_f16x2_avar,
                              NVPTX::ST_f32_avar, NVPTX::ST_f64_avar);
     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)
+  } 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,
                              NVPTX::ST_f32_asi, NVPTX::ST_f64_asi);
     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)
+  } else if (PointerSize == 64 ? SelectADDRri64(N2.getNode(), N2, Base, Offset)
                                : SelectADDRri(N2.getNode(), N2, Base, Offset)) {
-    if (TM.is64Bit())
+    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,
           NVPTX::ST_f16x2_ari_64, NVPTX::ST_f32_ari_64, NVPTX::ST_f64_ari_64);
     else
       Opcode = pickOpcodeForVT(SourceVT, NVPTX::ST_i8_ari, NVPTX::ST_i16_ari,
                                NVPTX::ST_i32_ari, NVPTX::ST_i64_ari,
                                NVPTX::ST_f16_ari, NVPTX::ST_f16x2_ari,
                                NVPTX::ST_f32_ari, NVPTX::ST_f64_ari);
     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())
+    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,
                           NVPTX::ST_f16_areg_64, NVPTX::ST_f16x2_areg_64,
                           NVPTX::ST_f32_areg_64, NVPTX::ST_f64_areg_64);
     else
       Opcode = pickOpcodeForVT(SourceVT, NVPTX::ST_i8_areg, NVPTX::ST_i16_areg,
                                NVPTX::ST_i32_areg, NVPTX::ST_i64_areg,
                                NVPTX::ST_f16_areg, NVPTX::ST_f16x2_areg,
                                NVPTX::ST_f32_areg, NVPTX::ST_f64_areg);
     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);
   }
 
   if (!NVPTXST)
     return false;
 
bool NVPTXDAGToDAGISel::tryStoreVector(SDNode *N) {
   SDLoc DL(N);
   SDNode *ST;
   EVT EltVT = Op1.getValueType();
   MemSDNode *MemSD = cast<MemSDNode>(N);
   EVT StoreVT = MemSD->getMemoryVT();
 
   // 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
   bool IsVolatile = MemSD->isVolatile();
   if (CodeAddrSpace != NVPTX::PTXLdStInstCode::GLOBAL &&
       CodeAddrSpace != NVPTX::PTXLdStInstCode::SHARED &&
       CodeAddrSpace != NVPTX::PTXLdStInstCode::GENERIC)
     IsVolatile = false;
case NVPTXISD::StoreV4:
       Opcode =
           pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy, NVPTX::STV_i8_v4_avar,
                           NVPTX::STV_i16_v4_avar, NVPTX::STV_i32_v4_avar, None,
                           NVPTX::STV_f16_v4_avar, NVPTX::STV_f16x2_v4_avar,
                           NVPTX::STV_f32_v4_avar, None);
       break;
     }
     StOps.push_back(Addr);
-  } else if (TM.is64Bit() ? SelectADDRsi64(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;
     case NVPTXISD::StoreV2:
       Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
                                NVPTX::STV_i8_v2_asi, NVPTX::STV_i16_v2_asi,
                                NVPTX::STV_i32_v2_asi, NVPTX::STV_i64_v2_asi,
                                NVPTX::STV_f16_v2_asi, NVPTX::STV_f16x2_v2_asi,
                                NVPTX::STV_f32_v2_asi, NVPTX::STV_f64_v2_asi);
       break;
     case NVPTXISD::StoreV4:
       Opcode =
           pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy, NVPTX::STV_i8_v4_asi,
                           NVPTX::STV_i16_v4_asi, NVPTX::STV_i32_v4_asi, None,
                           NVPTX::STV_f16_v4_asi, NVPTX::STV_f16x2_v4_asi,
                           NVPTX::STV_f32_v4_asi, None);
       break;
     }
     StOps.push_back(Base);
     StOps.push_back(Offset);
-  } else if (TM.is64Bit() ? SelectADDRri64(N2.getNode(), N2, Base, Offset)
+  } else if (PointerSize == 64 ? SelectADDRri64(N2.getNode(), N2, Base, Offset)
                                : SelectADDRri(N2.getNode(), N2, Base, Offset)) {
-    if (TM.is64Bit()) {
+    if (PointerSize == 64) {
       switch (N->getOpcode()) {
       default:
         return false;
       case NVPTXISD::StoreV2:
         Opcode = pickOpcodeForVT(
             EltVT.getSimpleVT().SimpleTy, NVPTX::STV_i8_v2_ari_64,
             NVPTX::STV_i16_v2_ari_64, NVPTX::STV_i32_v2_ari_64,
             NVPTX::STV_i64_v2_ari_64, NVPTX::STV_f16_v2_ari_64,
if (PointerSize == 64) {
                             NVPTX::STV_f16_v4_ari, NVPTX::STV_f16x2_v4_ari,
                             NVPTX::STV_f32_v4_ari, None);
         break;
       }
     }
     StOps.push_back(Base);
     StOps.push_back(Offset);
   } else {
-    if (TM.is64Bit()) {
+    if (PointerSize == 64) {
       switch (N->getOpcode()) {
       default:
         return false;
       case NVPTXISD::StoreV2:
         Opcode = pickOpcodeForVT(
             EltVT.getSimpleVT().SimpleTy, NVPTX::STV_i8_v2_areg_64,
             NVPTX::STV_i16_v2_areg_64, NVPTX::STV_i32_v2_areg_64,
             NVPTX::STV_i64_v2_areg_64, NVPTX::STV_f16_v2_areg_64,

llvm/trunk/lib/Target/NVPTX/NVPTXISelLowering.cpp

else {
           MakeIntrinsicCall(Intrinsic::nvvm_rsqrt_approx_d));
     }
   }
 }
 
 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);
 }
 
 std::string NVPTXTargetLowering::getPrototype(
     const DataLayout &DL, Type *retTy, const ArgListTy &Args,
     const SmallVectorImpl<ISD::OutputArg> &Outs, unsigned retAlignment,
     ImmutableCallSite CS) const {
   auto PtrVT = getPointerTy(DL);

llvm/trunk/lib/Target/NVPTX/NVPTXInstrInfo.td

 
 def hasPTX31 : Predicate<"Subtarget->getPTXVersion() >= 31">;
 def hasPTX60 : Predicate<"Subtarget->getPTXVersion() >= 60">;
 def hasPTX61 : Predicate<"Subtarget->getPTXVersion() >= 61">;
 
 def hasSM30 : Predicate<"Subtarget->getSmVersion() >= 30">;
 def hasSM70 : Predicate<"Subtarget->getSmVersion() >= 70">;
 
+def useShortPtr : Predicate<"useShortPointers()">;
 def useFP16Math: Predicate<"Subtarget->allowFP16Math()">;
 
 //===----------------------------------------------------------------------===//
 // Some Common Instruction Class Templates
 //===----------------------------------------------------------------------===//
 
 // Template for instructions which take three int64, int32, or int16 args.
 // The instructions are named "<OpcStr><Width>" (e.g. "add.s64").

llvm/trunk/lib/Target/NVPTX/NVPTXIntrinsics.td

 
 multiclass NG_TO_G<string Str, Intrinsic Intrin> {
    def _yes : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
           !strconcat("cvta.", Str, ".u32 \t$result, $src;"),
       [(set Int32Regs:$result, (Intrin Int32Regs:$src))]>;
    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> {
    def _yes : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
           !strconcat("cvta.to.", Str, ".u32 \t$result, $src;"),
       [(set Int32Regs:$result, (Intrin Int32Regs:$src))]>;
    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>;
 defm cvta_shared : NG_TO_G<"shared", int_nvvm_ptr_shared_to_gen>;
 defm cvta_global : NG_TO_G<"global", int_nvvm_ptr_global_to_gen>;
 defm cvta_const  : NG_TO_G<"const", int_nvvm_ptr_constant_to_gen>;
 
 defm cvta_to_local   : G_TO_NG<"local", int_nvvm_ptr_gen_to_local>;

llvm/trunk/lib/Target/NVPTX/NVPTXSubtarget.h

public:
   bool hasAtomBitwise64() const { return SmVersion >= 32; }
   bool hasAtomMinMax64() const { return SmVersion >= 32; }
   bool hasLDG() const { return SmVersion >= 32; }
   bool hasLDU() const { return ((SmVersion >= 20) && (SmVersion < 30)); }
   inline bool hasHWROT32() const { return SmVersion >= 32; }
   bool hasImageHandles() const;
   bool hasFP16Math() const { return SmVersion >= 53; }
   bool allowFP16Math() const;
-
   unsigned int getSmVersion() const { return SmVersion; }
   std::string getTargetName() const { return TargetName; }
 
   unsigned getPTXVersion() const { return PTXVersion; }
 
   NVPTXSubtarget &initializeSubtargetDependencies(StringRef CPU, StringRef FS);
   void ParseSubtargetFeatures(StringRef CPU, StringRef FS);
 };
 
 } // End llvm namespace
 
 #endif

llvm/trunk/lib/Target/NVPTX/NVPTXTargetMachine.h

 #include "llvm/Target/TargetMachine.h"
 
 namespace llvm {
 
 /// NVPTXTargetMachine
 ///
 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;
 
   // Hold Strings that can be free'd all together with NVPTXTargetMachine
   ManagedStringPool ManagedStrPool;
 
 public:
   NVPTXTargetMachine(const Target &T, const Triple &TT, StringRef CPU,
                      StringRef FS, const TargetOptions &Options,
                      Optional<Reloc::Model> RM, Optional<CodeModel::Model> CM,
                      CodeGenOpt::Level OP, bool is64bit);
 
   ~NVPTXTargetMachine() override;
   const NVPTXSubtarget *getSubtargetImpl(const Function &) const override {
     return &Subtarget;
   }
   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);
   }
 
   TargetPassConfig *createPassConfig(PassManagerBase &PM) override;
 
   // Emission of machine code through MCJIT is not supported.

llvm/trunk/lib/Target/NVPTX/NVPTXTargetMachine.cpp

 // TODO: Remove this flag when we are confident with no regressions.
 static cl::opt<bool> DisableRequireStructuredCFG(
     "disable-nvptx-require-structured-cfg",
     cl::desc("Transitional flag to turn off NVPTX's requirement on preserving "
              "structured CFG. The requirement should be disabled only when "
              "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&);
 void initializeNVVMReflectPass(PassRegistry&);
 void initializeGenericToNVVMPass(PassRegistry&);
 void initializeNVPTXAllocaHoistingPass(PassRegistry &);
 void initializeNVPTXAssignValidGlobalNamesPass(PassRegistry&);
 void initializeNVPTXLowerAggrCopiesPass(PassRegistry &);
extern "C" void LLVMInitializeNVPTXTarget() {
   initializeGenericToNVVMPass(PR);
   initializeNVPTXAllocaHoistingPass(PR);
   initializeNVPTXAssignValidGlobalNamesPass(PR);
   initializeNVPTXLowerArgsPass(PR);
   initializeNVPTXLowerAllocaPass(PR);
   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";
 
   return Ret;
 }
 
 static CodeModel::Model getEffectiveCodeModel(Optional<CodeModel::Model> CM) {
   if (CM)
     return *CM;
   return CodeModel::Small;
 }
 
 NVPTXTargetMachine::NVPTXTargetMachine(const Target &T, const Triple &TT,
                                        StringRef CPU, StringRef FS,
                                        const TargetOptions &Options,
                                        Optional<Reloc::Model> RM,
                                        Optional<CodeModel::Model> CM,
                                        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;
   else
     drvInterface = NVPTX::CUDA;
   if (!DisableRequireStructuredCFG)
     setRequiresStructuredCFG(true);
   initAsmInfo();

llvm/trunk/test/CodeGen/NVPTX/addrspacecast.ll

-; 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
 }

llvm/trunk/test/CodeGen/NVPTX/ld-addrspace.ll

-; 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
 }

llvm/trunk/test/CodeGen/NVPTX/st-addrspace.ll

-; 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
 }