Index: llvm/lib/Target/NVPTX/NVPTX.td =================================================================== --- llvm/lib/Target/NVPTX/NVPTX.td +++ llvm/lib/Target/NVPTX/NVPTX.td @@ -73,6 +73,11 @@ def PTX61 : SubtargetFeature<"ptx61", "PTXVersion", "61", "Use PTX version 6.1">; +// Other features +def SHORTPTR : SubtargetFeature<"short-ptr", "UseShortPointers", "true", + "Use 32-bit pointers for const/local/shared AS">; + + //===----------------------------------------------------------------------===// // NVPTX supported processors. //===----------------------------------------------------------------------===// Index: llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp =================================================================== --- llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp +++ llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp @@ -732,10 +732,10 @@ AddrSpaceCastSDNode *CastN = cast(N); unsigned SrcAddrSpace = CastN->getSrcAddressSpace(); unsigned DstAddrSpace = CastN->getDestAddressSpace(); - assert(SrcAddrSpace != DstAddrSpace && "addrspacecast must be between different address spaces"); + bool UseShortPointers = Subtarget->useShortPointers(); if (DstAddrSpace == ADDRESS_SPACE_GENERIC) { // Specific to generic unsigned Opc; @@ -745,13 +745,19 @@ 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 +775,19 @@ : 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 +843,23 @@ 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); } + bool UseShortPointers = Subtarget->useShortPointers() && + (CodeAddrSpace == NVPTX::PTXLdStInstCode::CONSTANT || + CodeAddrSpace == NVPTX::PTXLdStInstCode::LOCAL || + CodeAddrSpace == NVPTX::PTXLdStInstCode::SHARED); + unsigned int PointerSize = (!TM.is64Bit() || UseShortPointers) ? 32 : 64; + // 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,7 +906,7 @@ 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, @@ -905,14 +919,14 @@ 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()) + 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 @@ 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 @@ 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,15 @@ // Address Space Setting unsigned int CodeAddrSpace = getCodeAddrSpace(MemSD); - if (canLowerToLDG(MemSD, *Subtarget, CodeAddrSpace, MF)) { return tryLDGLDU(N); } + bool UseShortPointers = Subtarget->useShortPointers() && + (CodeAddrSpace == NVPTX::PTXLdStInstCode::CONSTANT || + CodeAddrSpace == NVPTX::PTXLdStInstCode::LOCAL || + CodeAddrSpace == NVPTX::PTXLdStInstCode::SHARED); + unsigned int PointerSize = (!TM.is64Bit() || UseShortPointers) ? 32 : 64; // Volatile Setting // - .volatile is only availalble for .global and .shared bool IsVolatile = MemSD->isVolatile(); @@ -1092,7 +1110,7 @@ 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()) { + if (PointerSize == 64) { switch (N->getOpcode()) { default: return false; @@ -1140,7 +1158,7 @@ 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,19 @@ return false; // Address Space Setting - unsigned int codeAddrSpace = getCodeAddrSpace(ST); + unsigned int CodeAddrSpace = getCodeAddrSpace(ST); + bool UseShortPointers = Subtarget->useShortPointers() && + (CodeAddrSpace == NVPTX::PTXLdStInstCode::CONSTANT || + CodeAddrSpace == NVPTX::PTXLdStInstCode::LOCAL || + CodeAddrSpace == NVPTX::PTXLdStInstCode::SHARED); + unsigned int PointerSize = (!TM.is64Bit() || UseShortPointers) ? 32 : 64; // 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,7 +1758,7 @@ 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); @@ -1748,13 +1771,13 @@ 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()) + 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 +1791,12 @@ 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 +1810,7 @@ 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,12 +1839,17 @@ // 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"); } + bool UseShortPointers = Subtarget->useShortPointers() && + (CodeAddrSpace == NVPTX::PTXLdStInstCode::CONSTANT || + CodeAddrSpace == NVPTX::PTXLdStInstCode::LOCAL || + CodeAddrSpace == NVPTX::PTXLdStInstCode::SHARED); + unsigned int PointerSize = (!TM.is64Bit() || UseShortPointers) ? 32 : 64; + // Volatile Setting // - .volatile is only availalble for .global and .shared bool IsVolatile = MemSD->isVolatile(); @@ -1925,7 +1953,7 @@ StOps.push_back(Offset); } else if (TM.is64Bit() ? SelectADDRri64(N2.getNode(), N2, Base, Offset) : SelectADDRri(N2.getNode(), N2, Base, Offset)) { - if (TM.is64Bit()) { + if (PointerSize == 64) { switch (N->getOpcode()) { default: return false; @@ -1968,7 +1996,7 @@ StOps.push_back(Base); StOps.push_back(Offset); } else { - if (TM.is64Bit()) { + if (PointerSize == 64) { switch (N->getOpcode()) { default: return false; Index: llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp =================================================================== --- llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp +++ llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp @@ -1243,9 +1243,18 @@ SDValue NVPTXTargetLowering::LowerGlobalAddress(SDValue Op, SelectionDAG &DAG) const { SDLoc dl(Op); - const GlobalValue *GV = cast(Op)->getGlobal(); - auto PtrVT = getPointerTy(DAG.getDataLayout()); - Op = DAG.getTargetGlobalAddress(GV, dl, PtrVT); + const GlobalAddressSDNode *GAN = cast(Op); + const NVPTXTargetMachine &TM = + static_cast(DAG.getTarget()); + + unsigned AS = GAN->getAddressSpace(); + bool UseShortPointer = TM.is64Bit() && STI.useShortPointers() && + (AS == ADDRESS_SPACE_CONST || ADDRESS_SPACE_LOCAL || + ADDRESS_SPACE_SHARED); + auto PtrVT = UseShortPointer + ? MVT::getIntegerVT(32) + : getPointerTy(DAG.getDataLayout(), GAN->getAddressSpace()); + Op = DAG.getTargetGlobalAddress(GAN->getGlobal(), dl, PtrVT); return DAG.getNode(NVPTXISD::Wrapper, dl, PtrVT, Op); } Index: llvm/lib/Target/NVPTX/NVPTXInstrInfo.td =================================================================== --- llvm/lib/Target/NVPTX/NVPTXInstrInfo.td +++ llvm/lib/Target/NVPTX/NVPTXInstrInfo.td @@ -147,6 +147,7 @@ def hasSM30 : Predicate<"Subtarget->getSmVersion() >= 30">; def hasSM70 : Predicate<"Subtarget->getSmVersion() >= 70">; +def useShortPtr : Predicate<"Subtarget->useShortPointers()">; def useFP16Math: Predicate<"Subtarget->allowFP16Math()">; //===----------------------------------------------------------------------===// Index: llvm/lib/Target/NVPTX/NVPTXIntrinsics.td =================================================================== --- llvm/lib/Target/NVPTX/NVPTXIntrinsics.td +++ llvm/lib/Target/NVPTX/NVPTXIntrinsics.td @@ -1937,6 +1937,12 @@ 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 { @@ -1946,6 +1952,12 @@ 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>; Index: llvm/lib/Target/NVPTX/NVPTXSubtarget.h =================================================================== --- llvm/lib/Target/NVPTX/NVPTXSubtarget.h +++ llvm/lib/Target/NVPTX/NVPTXSubtarget.h @@ -39,6 +39,9 @@ // SM version x.y is represented as 10*x+y, e.g. 3.1 == 31 unsigned int SmVersion; + // Use 32-bit pointers for accessing const/local/short AS. + bool UseShortPointers = false; + const NVPTXTargetMachine &TM; NVPTXInstrInfo InstrInfo; NVPTXTargetLowering TLInfo; @@ -79,7 +82,7 @@ bool hasImageHandles() const; bool hasFP16Math() const { return SmVersion >= 53; } bool allowFP16Math() const; - + bool useShortPointers() const { return UseShortPointers; } unsigned int getSmVersion() const { return SmVersion; } std::string getTargetName() const { return TargetName; } Index: llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp =================================================================== --- llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp +++ llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp @@ -83,11 +83,27 @@ initializeNVPTXLowerAggrCopiesPass(PR); } -static std::string computeDataLayout(bool is64Bit) { +static bool HasFeature(StringRef FS, StringRef WantedFeature) { + SmallVector Features; + FS.split(Features, ',', -1, false /* KeepEmpty */); + bool Result = false; + for (auto &Feature : Features) { + if (Feature == WantedFeature || + (Feature.startswith("+") && Feature.drop_front(1) == WantedFeature)) + Result = true; + else if (Feature.startswith("-") && Feature.drop_front(1) == WantedFeature) + Result = false; + } + return Result; +} + +static std::string computeDataLayout(bool is64Bit, StringRef FS) { std::string Ret = "e"; if (!is64Bit) Ret += "-p:32:32"; + else if (HasFeature(FS, "short-ptr")) + Ret += "-p3:32:32-p4:32:32-p5:32:32"; Ret += "-i64:64-i128:128-v16:16-v32:32-n16:32:64"; @@ -108,7 +124,7 @@ 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, + : LLVMTargetMachine(T, computeDataLayout(is64bit, FS), TT, CPU, FS, Options, Reloc::PIC_, getEffectiveCodeModel(CM), OL), is64bit(is64bit), TLOF(llvm::make_unique()), Subtarget(TT, CPU, FS, *this) { Index: llvm/test/CodeGen/NVPTX/addrspacecast.ll =================================================================== --- llvm/test/CodeGen/NVPTX/addrspacecast.ll +++ llvm/test/CodeGen/NVPTX/addrspacecast.ll @@ -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 -mattr=+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 Index: llvm/test/CodeGen/NVPTX/ld-addrspace.ll =================================================================== --- llvm/test/CodeGen/NVPTX/ld-addrspace.ll +++ llvm/test/CodeGen/NVPTX/ld-addrspace.ll @@ -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 -mattr=+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 } Index: llvm/test/CodeGen/NVPTX/st-addrspace.ll =================================================================== --- llvm/test/CodeGen/NVPTX/st-addrspace.ll +++ llvm/test/CodeGen/NVPTX/st-addrspace.ll @@ -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 -mattr=+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 }