[llvm] r185177 - [NVPTX] Calling conventions fix
Justin Holewinski
jholewinski at nvidia.com
Fri Jun 28 10:58:10 PDT 2013
Author: jholewinski
Date: Fri Jun 28 12:58:10 2013
New Revision: 185177
URL: http://llvm.org/viewvc/llvm-project?rev=185177&view=rev
Log:
[NVPTX] Calling conventions fix
Fix ABI handling for function
returning bool -- use st.param.b32 to return the value
and use ld.param.b32 in caller to load the return value.
Added:
llvm/trunk/test/CodeGen/NVPTX/i8-param.ll
Modified:
llvm/trunk/lib/Target/NVPTX/NVPTXISelLowering.cpp
llvm/trunk/test/CodeGen/NVPTX/compare-int.ll
llvm/trunk/test/CodeGen/NVPTX/convert-int-sm20.ll
llvm/trunk/test/CodeGen/NVPTX/ld-addrspace.ll
llvm/trunk/test/CodeGen/NVPTX/ld-generic.ll
Modified: llvm/trunk/lib/Target/NVPTX/NVPTXISelLowering.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXISelLowering.cpp?rev=185177&r1=185176&r2=185177&view=diff
==============================================================================
--- llvm/trunk/lib/Target/NVPTX/NVPTXISelLowering.cpp (original)
+++ llvm/trunk/lib/Target/NVPTX/NVPTXISelLowering.cpp Fri Jun 28 12:58:10 2013
@@ -1207,7 +1207,14 @@ SDValue NVPTXTargetLowering::LowerCall(T
sz = 8;
SmallVector<EVT, 4> LoadRetVTs;
- if (sz < 16) {
+ EVT TheLoadType = VTs[i];
+ if (retTy->isIntegerTy() &&
+ TD->getTypeAllocSizeInBits(retTy) < 32) {
+ // This is for integer types only, and specifically not for
+ // aggregates.
+ LoadRetVTs.push_back(MVT::i32);
+ TheLoadType = MVT::i32;
+ } else if (sz < 16) {
// If loading i1/i8 result, generate
// load i8 (-> i16)
// trunc i16 to i1/i8
@@ -1225,7 +1232,7 @@ SDValue NVPTXTargetLowering::LowerCall(T
SDValue retval = DAG.getMemIntrinsicNode(
NVPTXISD::LoadParam, dl,
DAG.getVTList(&LoadRetVTs[0], LoadRetVTs.size()), &LoadRetOps[0],
- LoadRetOps.size(), VTs[i], MachinePointerInfo());
+ LoadRetOps.size(), TheLoadType, MachinePointerInfo());
Chain = retval.getValue(1);
InFlag = retval.getValue(2);
SDValue Ret0 = retval.getValue(0);
@@ -1798,7 +1805,7 @@ NVPTXTargetLowering::LowerReturn(SDValue
SDLoc dl, SelectionDAG &DAG) const {
MachineFunction &MF = DAG.getMachineFunction();
const Function *F = MF.getFunction();
- const Type *RetTy = F->getReturnType();
+ Type *RetTy = F->getReturnType();
const DataLayout *TD = getDataLayout();
bool isABI = (nvptxSubtarget.getSmVersion() >= 20);
@@ -1806,14 +1813,14 @@ NVPTXTargetLowering::LowerReturn(SDValue
if (!isABI)
return Chain;
- if (const VectorType *VTy = dyn_cast<const VectorType>(RetTy)) {
+ if (VectorType *VTy = dyn_cast<VectorType>(RetTy)) {
// If we have a vector type, the OutVals array will be the scalarized
// components and we have combine them into 1 or more vector stores.
unsigned NumElts = VTy->getNumElements();
assert(NumElts == Outs.size() && "Bad scalarization of return value");
// const_cast can be removed in later LLVM versions
- EVT EltVT = getValueType(const_cast<Type *>(RetTy)).getVectorElementType();
+ EVT EltVT = getValueType(RetTy).getVectorElementType();
bool NeedExtend = false;
if (EltVT.getSizeInBits() < 16)
NeedExtend = true;
@@ -1923,34 +1930,43 @@ NVPTXTargetLowering::LowerReturn(SDValue
SmallVector<EVT, 16> ValVTs;
// const_cast is necessary since we are still using an LLVM version from
// before the type system re-write.
- ComputePTXValueVTs(*this, const_cast<Type *>(RetTy), ValVTs);
+ ComputePTXValueVTs(*this, RetTy, ValVTs);
assert(ValVTs.size() == OutVals.size() && "Bad return value decomposition");
- unsigned sizesofar = 0;
+ unsigned SizeSoFar = 0;
for (unsigned i = 0, e = Outs.size(); i != e; ++i) {
SDValue theVal = OutVals[i];
- EVT theValType = theVal.getValueType();
+ EVT TheValType = theVal.getValueType();
unsigned numElems = 1;
- if (theValType.isVector())
- numElems = theValType.getVectorNumElements();
+ if (TheValType.isVector())
+ numElems = TheValType.getVectorNumElements();
for (unsigned j = 0, je = numElems; j != je; ++j) {
- SDValue tmpval = theVal;
- if (theValType.isVector())
- tmpval = DAG.getNode(ISD::EXTRACT_VECTOR_ELT, dl,
- theValType.getVectorElementType(), tmpval,
+ SDValue TmpVal = theVal;
+ if (TheValType.isVector())
+ TmpVal = DAG.getNode(ISD::EXTRACT_VECTOR_ELT, dl,
+ TheValType.getVectorElementType(), TmpVal,
DAG.getIntPtrConstant(j));
- EVT theStoreType = tmpval.getValueType();
- if (theStoreType.getSizeInBits() < 8)
- tmpval = DAG.getNode(ISD::ZERO_EXTEND, dl, MVT::i8, tmpval);
- SDValue Ops[] = { Chain, DAG.getConstant(sizesofar, MVT::i32), tmpval };
+ EVT TheStoreType = ValVTs[i];
+ if (RetTy->isIntegerTy() &&
+ TD->getTypeAllocSizeInBits(RetTy) < 32) {
+ // The following zero-extension is for integer types only, and
+ // specifically not for aggregates.
+ TmpVal = DAG.getNode(ISD::ZERO_EXTEND, dl, MVT::i32, TmpVal);
+ TheStoreType = MVT::i32;
+ }
+ else if (TmpVal.getValueType().getSizeInBits() < 16)
+ TmpVal = DAG.getNode(ISD::ANY_EXTEND, dl, MVT::i16, TmpVal);
+
+ SDValue Ops[] = { Chain, DAG.getConstant(SizeSoFar, MVT::i32), TmpVal };
Chain = DAG.getMemIntrinsicNode(NVPTXISD::StoreRetval, dl,
- DAG.getVTList(MVT::Other), &Ops[0], 3,
- ValVTs[i], MachinePointerInfo());
- if (theValType.isVector())
- sizesofar +=
- ValVTs[i].getVectorElementType().getStoreSizeInBits() / 8;
+ DAG.getVTList(MVT::Other), &Ops[0],
+ 3, TheStoreType,
+ MachinePointerInfo());
+ if(TheValType.isVector())
+ SizeSoFar +=
+ TheStoreType.getVectorElementType().getStoreSizeInBits() / 8;
else
- sizesofar += ValVTs[i].getStoreSizeInBits() / 8;
+ SizeSoFar += TheStoreType.getStoreSizeInBits()/8;
}
}
}
Modified: llvm/trunk/test/CodeGen/NVPTX/compare-int.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/NVPTX/compare-int.ll?rev=185177&r1=185176&r2=185177&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/NVPTX/compare-int.ll (original)
+++ llvm/trunk/test/CodeGen/NVPTX/compare-int.ll Fri Jun 28 12:58:10 2013
@@ -195,7 +195,7 @@ define i32 @icmp_sle_i32(i32 %a, i32 %b)
define i16 @icmp_eq_i16(i16 %a, i16 %b) {
; CHECK: setp.eq.s16 %p[[P0:[0-9]+]], %rs{{[0-9]+}}, %rs{{[0-9]+}}
-; CHECK: selp.u16 %rs{{[0-9]+}}, 1, 0, %p[[P0]]
+; CHECK: selp.u32 %r{{[0-9]+}}, 1, 0, %p[[P0]]
; CHECK: ret
%cmp = icmp eq i16 %a, %b
%ret = zext i1 %cmp to i16
@@ -204,7 +204,7 @@ define i16 @icmp_eq_i16(i16 %a, i16 %b)
define i16 @icmp_ne_i16(i16 %a, i16 %b) {
; CHECK: setp.ne.s16 %p[[P0:[0-9]+]], %rs{{[0-9]+}}, %rs{{[0-9]+}}
-; CHECK: selp.u16 %rs{{[0-9]+}}, 1, 0, %p[[P0]]
+; CHECK: selp.u32 %r{{[0-9]+}}, 1, 0, %p[[P0]]
; CHECK: ret
%cmp = icmp ne i16 %a, %b
%ret = zext i1 %cmp to i16
@@ -213,7 +213,7 @@ define i16 @icmp_ne_i16(i16 %a, i16 %b)
define i16 @icmp_ugt_i16(i16 %a, i16 %b) {
; CHECK: setp.gt.u16 %p[[P0:[0-9]+]], %rs{{[0-9]+}}, %rs{{[0-9]+}}
-; CHECK: selp.u16 %rs{{[0-9]+}}, 1, 0, %p[[P0]]
+; CHECK: selp.u32 %r{{[0-9]+}}, 1, 0, %p[[P0]]
; CHECK: ret
%cmp = icmp ugt i16 %a, %b
%ret = zext i1 %cmp to i16
@@ -222,7 +222,7 @@ define i16 @icmp_ugt_i16(i16 %a, i16 %b)
define i16 @icmp_uge_i16(i16 %a, i16 %b) {
; CHECK: setp.ge.u16 %p[[P0:[0-9]+]], %rs{{[0-9]+}}, %rs{{[0-9]+}}
-; CHECK: selp.u16 %rs{{[0-9]+}}, 1, 0, %p[[P0]]
+; CHECK: selp.u32 %r{{[0-9]+}}, 1, 0, %p[[P0]]
; CHECK: ret
%cmp = icmp uge i16 %a, %b
%ret = zext i1 %cmp to i16
@@ -231,7 +231,7 @@ define i16 @icmp_uge_i16(i16 %a, i16 %b)
define i16 @icmp_ult_i16(i16 %a, i16 %b) {
; CHECK: setp.lt.u16 %p[[P0:[0-9]+]], %rs{{[0-9]+}}, %rs{{[0-9]+}}
-; CHECK: selp.u16 %rs{{[0-9]+}}, 1, 0, %p[[P0]]
+; CHECK: selp.u32 %r{{[0-9]+}}, 1, 0, %p[[P0]]
; CHECK: ret
%cmp = icmp ult i16 %a, %b
%ret = zext i1 %cmp to i16
@@ -240,7 +240,7 @@ define i16 @icmp_ult_i16(i16 %a, i16 %b)
define i16 @icmp_ule_i16(i16 %a, i16 %b) {
; CHECK: setp.le.u16 %p[[P0:[0-9]+]], %rs{{[0-9]+}}, %rs{{[0-9]+}}
-; CHECK: selp.u16 %rs{{[0-9]+}}, 1, 0, %p[[P0]]
+; CHECK: selp.u32 %r{{[0-9]+}}, 1, 0, %p[[P0]]
; CHECK: ret
%cmp = icmp ule i16 %a, %b
%ret = zext i1 %cmp to i16
@@ -249,7 +249,7 @@ define i16 @icmp_ule_i16(i16 %a, i16 %b)
define i16 @icmp_sgt_i16(i16 %a, i16 %b) {
; CHECK: setp.gt.s16 %p[[P0:[0-9]+]], %rs{{[0-9]+}}, %rs{{[0-9]+}}
-; CHECK: selp.u16 %rs{{[0-9]+}}, 1, 0, %p[[P0]]
+; CHECK: selp.u32 %r{{[0-9]+}}, 1, 0, %p[[P0]]
; CHECK: ret
%cmp = icmp sgt i16 %a, %b
%ret = zext i1 %cmp to i16
@@ -258,7 +258,7 @@ define i16 @icmp_sgt_i16(i16 %a, i16 %b)
define i16 @icmp_sge_i16(i16 %a, i16 %b) {
; CHECK: setp.ge.s16 %p[[P0:[0-9]+]], %rs{{[0-9]+}}, %rs{{[0-9]+}}
-; CHECK: selp.u16 %rs{{[0-9]+}}, 1, 0, %p[[P0]]
+; CHECK: selp.u32 %r{{[0-9]+}}, 1, 0, %p[[P0]]
; CHECK: ret
%cmp = icmp sge i16 %a, %b
%ret = zext i1 %cmp to i16
@@ -267,7 +267,7 @@ define i16 @icmp_sge_i16(i16 %a, i16 %b)
define i16 @icmp_slt_i16(i16 %a, i16 %b) {
; CHECK: setp.lt.s16 %p[[P0:[0-9]+]], %rs{{[0-9]+}}, %rs{{[0-9]+}}
-; CHECK: selp.u16 %rs{{[0-9]+}}, 1, 0, %p[[P0]]
+; CHECK: selp.u32 %r{{[0-9]+}}, 1, 0, %p[[P0]]
; CHECK: ret
%cmp = icmp slt i16 %a, %b
%ret = zext i1 %cmp to i16
@@ -276,7 +276,7 @@ define i16 @icmp_slt_i16(i16 %a, i16 %b)
define i16 @icmp_sle_i16(i16 %a, i16 %b) {
; CHECK: setp.le.s16 %p[[P0:[0-9]+]], %rs{{[0-9]+}}, %rs{{[0-9]+}}
-; CHECK: selp.u16 %rs{{[0-9]+}}, 1, 0, %p[[P0]]
+; CHECK: selp.u32 %r{{[0-9]+}}, 1, 0, %p[[P0]]
; CHECK: ret
%cmp = icmp sle i16 %a, %b
%ret = zext i1 %cmp to i16
@@ -289,7 +289,7 @@ define i16 @icmp_sle_i16(i16 %a, i16 %b)
define i8 @icmp_eq_i8(i8 %a, i8 %b) {
; Comparison happens in 16-bit
; CHECK: setp.eq.s16 %p[[P0:[0-9]+]], %rs{{[0-9]+}}, %rs{{[0-9]+}}
-; CHECK: selp.u16 %rs{{[0-9]+}}, 1, 0, %p[[P0]]
+; CHECK: selp.u32 %r{{[0-9]+}}, 1, 0, %p[[P0]]
; CHECK: ret
%cmp = icmp eq i8 %a, %b
%ret = zext i1 %cmp to i8
@@ -299,7 +299,7 @@ define i8 @icmp_eq_i8(i8 %a, i8 %b) {
define i8 @icmp_ne_i8(i8 %a, i8 %b) {
; Comparison happens in 16-bit
; CHECK: setp.ne.s16 %p[[P0:[0-9]+]], %rs{{[0-9]+}}, %rs{{[0-9]+}}
-; CHECK: selp.u16 %rs{{[0-9]+}}, 1, 0, %p[[P0]]
+; CHECK: selp.u32 %r{{[0-9]+}}, 1, 0, %p[[P0]]
; CHECK: ret
%cmp = icmp ne i8 %a, %b
%ret = zext i1 %cmp to i8
@@ -309,7 +309,7 @@ define i8 @icmp_ne_i8(i8 %a, i8 %b) {
define i8 @icmp_ugt_i8(i8 %a, i8 %b) {
; Comparison happens in 16-bit
; CHECK: setp.gt.u16 %p[[P0:[0-9]+]], %rs{{[0-9]+}}, %rs{{[0-9]+}}
-; CHECK: selp.u16 %rs{{[0-9]+}}, 1, 0, %p[[P0]]
+; CHECK: selp.u32 %r{{[0-9]+}}, 1, 0, %p[[P0]]
; CHECK: ret
%cmp = icmp ugt i8 %a, %b
%ret = zext i1 %cmp to i8
@@ -319,7 +319,7 @@ define i8 @icmp_ugt_i8(i8 %a, i8 %b) {
define i8 @icmp_uge_i8(i8 %a, i8 %b) {
; Comparison happens in 16-bit
; CHECK: setp.ge.u16 %p[[P0:[0-9]+]], %rs{{[0-9]+}}, %rs{{[0-9]+}}
-; CHECK: selp.u16 %rs{{[0-9]+}}, 1, 0, %p[[P0]]
+; CHECK: selp.u32 %r{{[0-9]+}}, 1, 0, %p[[P0]]
; CHECK: ret
%cmp = icmp uge i8 %a, %b
%ret = zext i1 %cmp to i8
@@ -329,7 +329,7 @@ define i8 @icmp_uge_i8(i8 %a, i8 %b) {
define i8 @icmp_ult_i8(i8 %a, i8 %b) {
; Comparison happens in 16-bit
; CHECK: setp.lt.u16 %p[[P0:[0-9]+]], %rs{{[0-9]+}}, %rs{{[0-9]+}}
-; CHECK: selp.u16 %rs{{[0-9]+}}, 1, 0, %p[[P0]]
+; CHECK: selp.u32 %r{{[0-9]+}}, 1, 0, %p[[P0]]
; CHECK: ret
%cmp = icmp ult i8 %a, %b
%ret = zext i1 %cmp to i8
@@ -339,7 +339,7 @@ define i8 @icmp_ult_i8(i8 %a, i8 %b) {
define i8 @icmp_ule_i8(i8 %a, i8 %b) {
; Comparison happens in 16-bit
; CHECK: setp.le.u16 %p[[P0:[0-9]+]], %rs{{[0-9]+}}, %rs{{[0-9]+}}
-; CHECK: selp.u16 %rs{{[0-9]+}}, 1, 0, %p[[P0]]
+; CHECK: selp.u32 %r{{[0-9]+}}, 1, 0, %p[[P0]]
; CHECK: ret
%cmp = icmp ule i8 %a, %b
%ret = zext i1 %cmp to i8
@@ -349,7 +349,7 @@ define i8 @icmp_ule_i8(i8 %a, i8 %b) {
define i8 @icmp_sgt_i8(i8 %a, i8 %b) {
; Comparison happens in 16-bit
; CHECK: setp.gt.s16 %p[[P0:[0-9]+]], %rs{{[0-9]+}}, %rs{{[0-9]+}}
-; CHECK: selp.u16 %rs{{[0-9]+}}, 1, 0, %p[[P0]]
+; CHECK: selp.u32 %r{{[0-9]+}}, 1, 0, %p[[P0]]
; CHECK: ret
%cmp = icmp sgt i8 %a, %b
%ret = zext i1 %cmp to i8
@@ -359,7 +359,7 @@ define i8 @icmp_sgt_i8(i8 %a, i8 %b) {
define i8 @icmp_sge_i8(i8 %a, i8 %b) {
; Comparison happens in 16-bit
; CHECK: setp.ge.s16 %p[[P0:[0-9]+]], %rs{{[0-9]+}}, %rs{{[0-9]+}}
-; CHECK: selp.u16 %rs{{[0-9]+}}, 1, 0, %p[[P0]]
+; CHECK: selp.u32 %r{{[0-9]+}}, 1, 0, %p[[P0]]
; CHECK: ret
%cmp = icmp sge i8 %a, %b
%ret = zext i1 %cmp to i8
@@ -369,7 +369,7 @@ define i8 @icmp_sge_i8(i8 %a, i8 %b) {
define i8 @icmp_slt_i8(i8 %a, i8 %b) {
; Comparison happens in 16-bit
; CHECK: setp.lt.s16 %p[[P0:[0-9]+]], %rs{{[0-9]+}}, %rs{{[0-9]+}}
-; CHECK: selp.u16 %rs{{[0-9]+}}, 1, 0, %p[[P0]]
+; CHECK: selp.u32 %r{{[0-9]+}}, 1, 0, %p[[P0]]
; CHECK: ret
%cmp = icmp slt i8 %a, %b
%ret = zext i1 %cmp to i8
@@ -379,7 +379,7 @@ define i8 @icmp_slt_i8(i8 %a, i8 %b) {
define i8 @icmp_sle_i8(i8 %a, i8 %b) {
; Comparison happens in 16-bit
; CHECK: setp.le.s16 %p[[P0:[0-9]+]], %rs{{[0-9]+}}, %rs{{[0-9]+}}
-; CHECK: selp.u16 %rs{{[0-9]+}}, 1, 0, %p[[P0]]
+; CHECK: selp.u32 %r{{[0-9]+}}, 1, 0, %p[[P0]]
; CHECK: ret
%cmp = icmp sle i8 %a, %b
%ret = zext i1 %cmp to i8
Modified: llvm/trunk/test/CodeGen/NVPTX/convert-int-sm20.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/NVPTX/convert-int-sm20.ll?rev=185177&r1=185176&r2=185177&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/NVPTX/convert-int-sm20.ll (original)
+++ llvm/trunk/test/CodeGen/NVPTX/convert-int-sm20.ll Fri Jun 28 12:58:10 2013
@@ -8,16 +8,16 @@
; i16
define i16 @cvt_i16_i32(i32 %x) {
-; CHECK: ld.param.u16 %rs[[R0:[0-9]+]], [cvt_i16_i32_param_{{[0-9]+}}]
-; CHECK: st.param.b16 [func_retval{{[0-9]+}}+0], %rs[[R0]]
+; CHECK: ld.param.u16 %r[[R0:[0-9]+]], [cvt_i16_i32_param_{{[0-9]+}}]
+; CHECK: st.param.b32 [func_retval{{[0-9]+}}+0], %r[[R0]]
; CHECK: ret
%a = trunc i32 %x to i16
ret i16 %a
}
define i16 @cvt_i16_i64(i64 %x) {
-; CHECK: ld.param.u16 %rs[[R0:[0-9]+]], [cvt_i16_i64_param_{{[0-9]+}}]
-; CHECK: st.param.b16 [func_retval{{[0-9]+}}+0], %rs[[R0]]
+; CHECK: ld.param.u16 %r[[R0:[0-9]+]], [cvt_i16_i64_param_{{[0-9]+}}]
+; CHECK: st.param.b32 [func_retval{{[0-9]+}}+0], %r[[R0]]
; CHECK: ret
%a = trunc i64 %x to i16
ret i16 %a
Added: llvm/trunk/test/CodeGen/NVPTX/i8-param.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/NVPTX/i8-param.ll?rev=185177&view=auto
==============================================================================
--- llvm/trunk/test/CodeGen/NVPTX/i8-param.ll (added)
+++ llvm/trunk/test/CodeGen/NVPTX/i8-param.ll Fri Jun 28 12:58:10 2013
@@ -0,0 +1,23 @@
+; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s
+
+target datalayout = "e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
+
+; CHECK: .visible .func (.param .b32 func_retval0) callee
+define i8 @callee(i8 %a) {
+; CHECK: ld.param.u8
+ %ret = add i8 %a, 42
+; CHECK: st.param.b32
+ ret i8 %ret
+}
+
+; CHECK: .visible .func caller
+define void @caller(i8* %a) {
+; CHECK: ld.u8
+ %val = load i8* %a
+ %ret = tail call i8 @callee(i8 %val)
+; CHECK: ld.param.b32
+ store i8 %ret, i8* %a
+ ret void
+}
+
+
\ No newline at end of file
Modified: llvm/trunk/test/CodeGen/NVPTX/ld-addrspace.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/NVPTX/ld-addrspace.ll?rev=185177&r1=185176&r2=185177&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/NVPTX/ld-addrspace.ll (original)
+++ llvm/trunk/test/CodeGen/NVPTX/ld-addrspace.ll Fri Jun 28 12:58:10 2013
@@ -4,27 +4,27 @@
;; i8
define i8 @ld_global_i8(i8 addrspace(1)* %ptr) {
-; PTX32: ld.global.u8 %rs{{[0-9]+}}, [%r{{[0-9]+}}]
+; PTX32: ld.global.u8 %r{{[0-9]+}}, [%r{{[0-9]+}}]
; PTX32: ret
-; PTX64: ld.global.u8 %rs{{[0-9]+}}, [%rl{{[0-9]+}}]
+; PTX64: ld.global.u8 %r{{[0-9]+}}, [%rl{{[0-9]+}}]
; PTX64: ret
%a = load i8 addrspace(1)* %ptr
ret i8 %a
}
define i8 @ld_shared_i8(i8 addrspace(3)* %ptr) {
-; PTX32: ld.shared.u8 %rs{{[0-9]+}}, [%r{{[0-9]+}}]
+; PTX32: ld.shared.u8 %r{{[0-9]+}}, [%r{{[0-9]+}}]
; PTX32: ret
-; PTX64: ld.shared.u8 %rs{{[0-9]+}}, [%rl{{[0-9]+}}]
+; PTX64: ld.shared.u8 %r{{[0-9]+}}, [%rl{{[0-9]+}}]
; PTX64: ret
%a = load i8 addrspace(3)* %ptr
ret i8 %a
}
define i8 @ld_local_i8(i8 addrspace(5)* %ptr) {
-; PTX32: ld.local.u8 %rs{{[0-9]+}}, [%r{{[0-9]+}}]
+; PTX32: ld.local.u8 %r{{[0-9]+}}, [%r{{[0-9]+}}]
; PTX32: ret
-; PTX64: ld.local.u8 %rs{{[0-9]+}}, [%rl{{[0-9]+}}]
+; PTX64: ld.local.u8 %r{{[0-9]+}}, [%rl{{[0-9]+}}]
; PTX64: ret
%a = load i8 addrspace(5)* %ptr
ret i8 %a
@@ -32,27 +32,27 @@ define i8 @ld_local_i8(i8 addrspace(5)*
;; i16
define i16 @ld_global_i16(i16 addrspace(1)* %ptr) {
-; PTX32: ld.global.u16 %rs{{[0-9]+}}, [%r{{[0-9]+}}]
+; PTX32: ld.global.u16 %r{{[0-9]+}}, [%r{{[0-9]+}}]
; PTX32: ret
-; PTX64: ld.global.u16 %rs{{[0-9]+}}, [%rl{{[0-9]+}}]
+; PTX64: ld.global.u16 %r{{[0-9]+}}, [%rl{{[0-9]+}}]
; PTX64: ret
%a = load i16 addrspace(1)* %ptr
ret i16 %a
}
define i16 @ld_shared_i16(i16 addrspace(3)* %ptr) {
-; PTX32: ld.shared.u16 %rs{{[0-9]+}}, [%r{{[0-9]+}}]
+; PTX32: ld.shared.u16 %r{{[0-9]+}}, [%r{{[0-9]+}}]
; PTX32: ret
-; PTX64: ld.shared.u16 %rs{{[0-9]+}}, [%rl{{[0-9]+}}]
+; PTX64: ld.shared.u16 %r{{[0-9]+}}, [%rl{{[0-9]+}}]
; PTX64: ret
%a = load i16 addrspace(3)* %ptr
ret i16 %a
}
define i16 @ld_local_i16(i16 addrspace(5)* %ptr) {
-; PTX32: ld.local.u16 %rs{{[0-9]+}}, [%r{{[0-9]+}}]
+; PTX32: ld.local.u16 %r{{[0-9]+}}, [%r{{[0-9]+}}]
; PTX32: ret
-; PTX64: ld.local.u16 %rs{{[0-9]+}}, [%rl{{[0-9]+}}]
+; PTX64: ld.local.u16 %r{{[0-9]+}}, [%rl{{[0-9]+}}]
; PTX64: ret
%a = load i16 addrspace(5)* %ptr
ret i16 %a
Modified: llvm/trunk/test/CodeGen/NVPTX/ld-generic.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/NVPTX/ld-generic.ll?rev=185177&r1=185176&r2=185177&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/NVPTX/ld-generic.ll (original)
+++ llvm/trunk/test/CodeGen/NVPTX/ld-generic.ll Fri Jun 28 12:58:10 2013
@@ -4,9 +4,9 @@
;; i8
define i8 @ld_global_i8(i8 addrspace(0)* %ptr) {
-; PTX32: ld.u8 %rs{{[0-9]+}}, [%r{{[0-9]+}}]
+; PTX32: ld.u8 %r{{[0-9]+}}, [%r{{[0-9]+}}]
; PTX32: ret
-; PTX64: ld.u8 %rs{{[0-9]+}}, [%rl{{[0-9]+}}]
+; PTX64: ld.u8 %r{{[0-9]+}}, [%rl{{[0-9]+}}]
; PTX64: ret
%a = load i8 addrspace(0)* %ptr
ret i8 %a
@@ -14,9 +14,9 @@ define i8 @ld_global_i8(i8 addrspace(0)*
;; i16
define i16 @ld_global_i16(i16 addrspace(0)* %ptr) {
-; PTX32: ld.u16 %rs{{[0-9]+}}, [%r{{[0-9]+}}]
+; PTX32: ld.u16 %r{{[0-9]+}}, [%r{{[0-9]+}}]
; PTX32: ret
-; PTX64: ld.u16 %rs{{[0-9]+}}, [%rl{{[0-9]+}}]
+; PTX64: ld.u16 %r{{[0-9]+}}, [%rl{{[0-9]+}}]
; PTX64: ret
%a = load i16 addrspace(0)* %ptr
ret i16 %a
More information about the llvm-commits
mailing list