Hi Justin, <br><br>The assertion on NVPTXISelLowering.cpp:3305<div><br></div><div>assert(I.hasMetadataOtherThanDebugLoc() && "Must have alignment metadata");</div><div><br></div><div>looks a little fragile. !align is extended metadata and the optimization passes (e.g., SimplifyCFG) don't guarantee to preserve it. </div>
<div><br></div><div>I attached a reduced test case (hoist.ll) to demonstrate this issue. If you run "opt -simplifycfg hoist.ll -S", you will see the !align metadata is removed by HoistThenElseCodeToIf. Further running llc on the simplified bitcode will hit the assertion error. </div>
<div><br></div><div>One way to fix this issue is of course having HositThenElseCodeToIf to preserve the align metadata. But before I do that, I am curious whether we should assume the align metadata always exists at first hand. Can you shed some light on this, Justin? </div>
<div><br></div><div>Thanks much,</div><div>Jingyue</div><br><div class="gmail_quote">On Fri Jun 27 2014 at 12:07:55 PM Justin Holewinski <<a href="mailto:jholewinski@nvidia.com" target="_blank">jholewinski@nvidia.com</a>> wrote:<br>
<blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">Author: jholewinski<br>
Date: Fri Jun 27 13:35:51 2014<br>
New Revision: 211939<br>
<br>
URL: <a href="http://llvm.org/viewvc/llvm-project?rev=211939&view=rev" target="_blank">http://llvm.org/viewvc/llvm-<u></u>pr<u></u>oject?rev=211939&view=rev</a><br>
Log:<br>
[NVPTX] Fix handling of ldg/ldu intrinsics.<br>
<br>
The address space of the pointer must be global (1) for these intrinsics. There must also be alignment metadata attached to the intrinsic calls, e.g.<br>
<br>
%val = tail call i32 @llvm.nvvm.ldu.i.global.i32.<u></u>p1<u></u>i32(i32 addrspace(1)* %ptr), !align !0<br>
<br>
!0 = metadata !{i32 4}<br>
<br>
Added:<br>
llvm/trunk/test/CodeGen/NVPTX/<u></u><u></u>ldu-ldg.ll<br>
Modified:<br>
llvm/trunk/include/llvm/IR/<u></u>Int<u></u>rinsicsNVVM.td<br>
llvm/trunk/lib/Target/NVPTX/<u></u>NV<u></u>PTXISelDAGToDAG.cpp<br>
llvm/trunk/lib/Target/NVPTX/<u></u>NV<u></u>PTXISelDAGToDAG.h<br>
llvm/trunk/lib/Target/NVPTX/<u></u>NV<u></u>PTXISelLowering.cpp<br>
llvm/trunk/lib/Target/NVPTX/<u></u>NV<u></u>PTXIntrinsics.td<br>
llvm/trunk/test/CodeGen/NVPTX/<u></u><u></u>ldu-i8.ll<br>
llvm/trunk/test/CodeGen/NVPTX/<u></u><u></u>ldu-reg-plus-offset.ll<br>
<br>
Modified: llvm/trunk/include/llvm/IR/<u></u>Int<u></u>rinsicsNVVM.td<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/IR/IntrinsicsNVVM.td?rev=211939&r1=211938&r2=211939&view=diff" target="_blank">http://llvm.org/viewvc/llvm-<u></u>pr<u></u>oject/llvm/trunk/include/<u></u>llvm/<u></u>IR/IntrinsicsNVVM.td?rev=<u></u>21193<u></u>9&r1=211938&r2=211939&<u></u>view=<u></u>diff</a><br>
==============================<u></u><u></u>==============================<u></u><u></u>==================<br>
--- llvm/trunk/include/llvm/IR/<u></u>Int<u></u>rinsicsNVVM.td (original)<br>
+++ llvm/trunk/include/llvm/IR/<u></u>Int<u></u>rinsicsNVVM.td Fri Jun 27 13:35:51 2014<br>
@@ -796,26 +796,25 @@ def llvm_anyi64ptr_ty : LLVMAnyPoint<br>
<br>
<br>
// Generated within nvvm. Use for ldu on sm_20 or later<br>
-// @TODO: Revisit this, Changed LLVMAnyPointerType to LLVMPointerType<br>
def int_nvvm_ldu_global_i : Intrinsic<[llvm_anyint_ty],<br>
- [LLVMPointerType<<u></u>LLVMMatchType<u></u><0>>], [IntrReadMem, NoCapture<0>],<br>
+ [LLVMAnyPointerType<<u></u>LLVMMatchT<u></u>ype<0>>], [IntrReadMem, NoCapture<0>],<br>
"llvm.nvvm.ldu.global.i">;<br>
def int_nvvm_ldu_global_f : Intrinsic<[llvm_anyfloat_ty],<br>
- [LLVMPointerType<<u></u>LLVMMatchType<u></u><0>>], [IntrReadMem, NoCapture<0>],<br>
+ [LLVMAnyPointerType<<u></u>LLVMMatchT<u></u>ype<0>>], [IntrReadMem, NoCapture<0>],<br>
"llvm.nvvm.ldu.global.f">;<br>
def int_nvvm_ldu_global_p : Intrinsic<[llvm_anyptr_ty],<br>
- [LLVMPointerType<<u></u>LLVMMatchType<u></u><0>>], [IntrReadMem, NoCapture<0>],<br>
+ [LLVMAnyPointerType<<u></u>LLVMMatchT<u></u>ype<0>>], [IntrReadMem, NoCapture<0>],<br>
"llvm.nvvm.ldu.global.p">;<br>
<br>
// Generated within nvvm. Use for ldg on sm_35 or later<br>
def int_nvvm_ldg_global_i : Intrinsic<[llvm_anyint_ty],<br>
- [LLVMPointerType<<u></u>LLVMMatchType<u></u><0>>], [IntrReadMem, NoCapture<0>],<br>
+ [LLVMAnyPointerType<<u></u>LLVMMatchT<u></u>ype<0>>], [IntrReadMem, NoCapture<0>],<br>
"llvm.nvvm.ldg.global.i">;<br>
def int_nvvm_ldg_global_f : Intrinsic<[llvm_anyfloat_ty],<br>
- [LLVMPointerType<<u></u>LLVMMatchType<u></u><0>>], [IntrReadMem, NoCapture<0>],<br>
+ [LLVMAnyPointerType<<u></u>LLVMMatchT<u></u>ype<0>>], [IntrReadMem, NoCapture<0>],<br>
"llvm.nvvm.ldg.global.f">;<br>
def int_nvvm_ldg_global_p : Intrinsic<[llvm_anyptr_ty],<br>
- [LLVMPointerType<<u></u>LLVMMatchType<u></u><0>>], [IntrReadMem, NoCapture<0>],<br>
+ [LLVMAnyPointerType<<u></u>LLVMMatchT<u></u>ype<0>>], [IntrReadMem, NoCapture<0>],<br>
"llvm.nvvm.ldg.global.p">;<br>
<br>
// Use for generic pointers<br>
<br>
Modified: llvm/trunk/lib/Target/NVPTX/<u></u>NV<u></u>PTXISelDAGToDAG.cpp<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp?rev=211939&r1=211938&r2=211939&view=diff" target="_blank">http://llvm.org/viewvc/llvm-<u></u>pr<u></u>oject/llvm/trunk/lib/Target/<u></u>NV<u></u>PTX/NVPTXISelDAGToDAG.cpp?<u></u>rev=<u></u>211939&r1=211938&r2=<u></u>211939&<u></u>view=diff</a><br>
==============================<u></u><u></u>==============================<u></u><u></u>==================<br>
--- llvm/trunk/lib/Target/NVPTX/<u></u>NV<u></u>PTXISelDAGToDAG.cpp (original)<br>
+++ llvm/trunk/lib/Target/NVPTX/<u></u>NV<u></u>PTXISelDAGToDAG.cpp Fri Jun 27 13:35:51 2014<br>
@@ -141,7 +141,7 @@ SDNode *NVPTXDAGToDAGISel::Select(<u></u>SDN<u></u>ode<br>
case NVPTXISD::LDGV4:<br>
case NVPTXISD::LDUV2:<br>
case NVPTXISD::LDUV4:<br>
- ResNode = SelectLDGLDUVector(N);<br>
+ ResNode = SelectLDGLDU(N);<br>
break;<br>
case NVPTXISD::StoreV2:<br>
case NVPTXISD::StoreV4:<br>
@@ -167,6 +167,9 @@ SDNode *NVPTXDAGToDAGISel::Select(<u></u>SDN<u></u>ode<br>
case ISD::INTRINSIC_WO_CHAIN:<br>
ResNode = SelectIntrinsicNoChain(N);<br>
break;<br>
+ case ISD::INTRINSIC_W_CHAIN:<br>
+ ResNode = SelectIntrinsicChain(N);<br>
+ break;<br>
case NVPTXISD::Tex1DFloatI32:<br>
case NVPTXISD::Tex1DFloatFloat:<br>
case NVPTXISD::<u></u>Tex1DFloatFloatLevel<u></u>:<br>
@@ -273,6 +276,21 @@ SDNode *NVPTXDAGToDAGISel::Select(<u></u>SDN<u></u>ode<br>
return SelectCode(N);<br>
}<br>
<br>
+SDNode *NVPTXDAGToDAGISel::<u></u>SelectIntr<u></u>insicChain(SDNode *N) {<br>
+ unsigned IID = cast<ConstantSDNode>(N-><u></u>getOpe<u></u>rand(1))->getZExtValue()<u></u>;<br>
+ switch (IID) {<br>
+ default:<br>
+ return NULL;<br>
+ case Intrinsic::nvvm_ldg_global_f:<br>
+ case Intrinsic::nvvm_ldg_global_i:<br>
+ case Intrinsic::nvvm_ldg_global_p:<br>
+ case Intrinsic::nvvm_ldu_global_f:<br>
+ case Intrinsic::nvvm_ldu_global_i:<br>
+ case Intrinsic::nvvm_ldu_global_p:<br>
+ return SelectLDGLDU(N);<br>
+ }<br>
+}<br>
+<br>
static unsigned int getCodeAddrSpace(MemSDNode *N,<br>
const NVPTXSubtarget &Subtarget) {<br>
const Value *Src = N->getMemOperand()->getValue()<u></u><u></u>;<br>
@@ -990,22 +1008,101 @@ SDNode *NVPTXDAGToDAGISel::<u></u>SelectLoad<u></u>Vec<br>
return LD;<br>
}<br>
<br>
-SDNode *NVPTXDAGToDAGISel::<u></u>SelectLDGL<u></u>DUVector(SDNode *N) {<br>
+SDNode *NVPTXDAGToDAGISel::<u></u>SelectLDGL<u></u>DU(SDNode *N) {<br>
<br>
SDValue Chain = N->getOperand(0);<br>
- SDValue Op1 = N->getOperand(1);<br>
+ SDValue Op1;<br>
+ MemSDNode *Mem;<br>
+ bool IsLDG = true;<br>
+<br>
+ // If this is an LDG intrinsic, the address is the third operand. Its its an<br>
+ // LDG/LDU SD node (from custom vector handling), then its the second operand<br>
+ if (N->getOpcode() == ISD::INTRINSIC_W_CHAIN) {<br>
+ Op1 = N->getOperand(2);<br>
+ Mem = cast<MemIntrinsicSDNode>(N);<br>
+ unsigned IID = cast<ConstantSDNode>(N-><u></u>getOpe<u></u>rand(1))->getZExtValue()<u></u>;<br>
+ switch (IID) {<br>
+ default:<br>
+ return NULL;<br>
+ case Intrinsic::nvvm_ldg_global_f:<br>
+ case Intrinsic::nvvm_ldg_global_i:<br>
+ case Intrinsic::nvvm_ldg_global_p:<br>
+ IsLDG = true;<br>
+ break;<br>
+ case Intrinsic::nvvm_ldu_global_f:<br>
+ case Intrinsic::nvvm_ldu_global_i:<br>
+ case Intrinsic::nvvm_ldu_global_p:<br>
+ IsLDG = false;<br>
+ break;<br>
+ }<br>
+ } else {<br>
+ Op1 = N->getOperand(1);<br>
+ Mem = cast<MemSDNode>(N);<br>
+ }<br>
+<br>
unsigned Opcode;<br>
SDLoc DL(N);<br>
SDNode *LD;<br>
- MemSDNode *Mem = cast<MemSDNode>(N);<br>
SDValue Base, Offset, Addr;<br>
<br>
- EVT EltVT = Mem->getMemoryVT().<u></u>getVectorEl<u></u>ementType();<br>
+ EVT EltVT = Mem->getMemoryVT();<br>
+ if (EltVT.isVector()) {<br>
+ EltVT = EltVT.getVectorElementType();<br>
+ }<br>
<br>
if (SelectDirectAddr(Op1, Addr)) {<br>
switch (N->getOpcode()) {<br>
default:<br>
return nullptr;<br>
+ case ISD::INTRINSIC_W_CHAIN:<br>
+ if (IsLDG) {<br>
+ switch (EltVT.getSimpleVT().SimpleTy) {<br>
+ default:<br>
+ return nullptr;<br>
+ case MVT::i8:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i8av<u></u>ar;<br>
+ break;<br>
+ case MVT::i16:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i16a<u></u>var;<br>
+ break;<br>
+ case MVT::i32:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i32a<u></u>var;<br>
+ break;<br>
+ case MVT::i64:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i64a<u></u>var;<br>
+ break;<br>
+ case MVT::f32:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>f32a<u></u>var;<br>
+ break;<br>
+ case MVT::f64:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>f64a<u></u>var;<br>
+ break;<br>
+ }<br>
+ } else {<br>
+ switch (EltVT.getSimpleVT().SimpleTy) {<br>
+ default:<br>
+ return nullptr;<br>
+ case MVT::i8:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i8av<u></u>ar;<br>
+ break;<br>
+ case MVT::i16:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i16a<u></u>var;<br>
+ break;<br>
+ case MVT::i32:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i32a<u></u>var;<br>
+ break;<br>
+ case MVT::i64:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i64a<u></u>var;<br>
+ break;<br>
+ case MVT::f32:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>f32a<u></u>var;<br>
+ break;<br>
+ case MVT::f64:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>f64a<u></u>var;<br>
+ break;<br>
+ }<br>
+ }<br>
+ break;<br>
case NVPTXISD::LDGV2:<br>
switch (EltVT.getSimpleVT().SimpleTy) {<br>
default:<br>
@@ -1101,6 +1198,55 @@ SDNode *NVPTXDAGToDAGISel::<u></u>SelectLDGL<u></u>DUV<br>
switch (N->getOpcode()) {<br>
default:<br>
return nullptr;<br>
+ case ISD::INTRINSIC_W_CHAIN:<br>
+ if (IsLDG) {<br>
+ switch (EltVT.getSimpleVT().SimpleTy) {<br>
+ default:<br>
+ return nullptr;<br>
+ case MVT::i8:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i8ar<u></u>i64;<br>
+ break;<br>
+ case MVT::i16:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i16a<u></u>ri64;<br>
+ break;<br>
+ case MVT::i32:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i32a<u></u>ri64;<br>
+ break;<br>
+ case MVT::i64:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i64a<u></u>ri64;<br>
+ break;<br>
+ case MVT::f32:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>f32a<u></u>ri64;<br>
+ break;<br>
+ case MVT::f64:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>f64a<u></u>ri64;<br>
+ break;<br>
+ }<br>
+ } else {<br>
+ switch (EltVT.getSimpleVT().SimpleTy) {<br>
+ default:<br>
+ return nullptr;<br>
+ case MVT::i8:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i8ar<u></u>i64;<br>
+ break;<br>
+ case MVT::i16:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i16a<u></u>ri64;<br>
+ break;<br>
+ case MVT::i32:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i32a<u></u>ri64;<br>
+ break;<br>
+ case MVT::i64:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i64a<u></u>ri64;<br>
+ break;<br>
+ case MVT::f32:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>f32a<u></u>ri64;<br>
+ break;<br>
+ case MVT::f64:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>f64a<u></u>ri64;<br>
+ break;<br>
+ }<br>
+ }<br>
+ break;<br>
case NVPTXISD::LDGV2:<br>
switch (EltVT.getSimpleVT().SimpleTy) {<br>
default:<br>
@@ -1190,6 +1336,55 @@ SDNode *NVPTXDAGToDAGISel::<u></u>SelectLDGL<u></u>DUV<br>
switch (N->getOpcode()) {<br>
default:<br>
return nullptr;<br>
+ case ISD::INTRINSIC_W_CHAIN:<br>
+ if (IsLDG) {<br>
+ switch (EltVT.getSimpleVT().SimpleTy) {<br>
+ default:<br>
+ return nullptr;<br>
+ case MVT::i8:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i8ar<u></u>i;<br>
+ break;<br>
+ case MVT::i16:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i16a<u></u>ri;<br>
+ break;<br>
+ case MVT::i32:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i32a<u></u>ri;<br>
+ break;<br>
+ case MVT::i64:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i64a<u></u>ri;<br>
+ break;<br>
+ case MVT::f32:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>f32a<u></u>ri;<br>
+ break;<br>
+ case MVT::f64:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>f64a<u></u>ri;<br>
+ break;<br>
+ }<br>
+ } else {<br>
+ switch (EltVT.getSimpleVT().SimpleTy) {<br>
+ default:<br>
+ return nullptr;<br>
+ case MVT::i8:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i8ar<u></u>i;<br>
+ break;<br>
+ case MVT::i16:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i16a<u></u>ri;<br>
+ break;<br>
+ case MVT::i32:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i32a<u></u>ri;<br>
+ break;<br>
+ case MVT::i64:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i64a<u></u>ri;<br>
+ break;<br>
+ case MVT::f32:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>f32a<u></u>ri;<br>
+ break;<br>
+ case MVT::f64:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>f64a<u></u>ri;<br>
+ break;<br>
+ }<br>
+ }<br>
+ break;<br>
case NVPTXISD::LDGV2:<br>
switch (EltVT.getSimpleVT().SimpleTy) {<br>
default:<br>
@@ -1285,6 +1480,55 @@ SDNode *NVPTXDAGToDAGISel::<u></u>SelectLDGL<u></u>DUV<br>
switch (N->getOpcode()) {<br>
default:<br>
return nullptr;<br>
+ case ISD::INTRINSIC_W_CHAIN:<br>
+ if (IsLDG) {<br>
+ switch (EltVT.getSimpleVT().SimpleTy) {<br>
+ default:<br>
+ return nullptr;<br>
+ case MVT::i8:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i8ar<u></u>eg64;<br>
+ break;<br>
+ case MVT::i16:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i16a<u></u>reg64;<br>
+ break;<br>
+ case MVT::i32:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i32a<u></u>reg64;<br>
+ break;<br>
+ case MVT::i64:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i64a<u></u>reg64;<br>
+ break;<br>
+ case MVT::f32:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>f32a<u></u>reg64;<br>
+ break;<br>
+ case MVT::f64:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>f64a<u></u>reg64;<br>
+ break;<br>
+ }<br>
+ } else {<br>
+ switch (EltVT.getSimpleVT().SimpleTy) {<br>
+ default:<br>
+ return nullptr;<br>
+ case MVT::i8:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i8ar<u></u>eg64;<br>
+ break;<br>
+ case MVT::i16:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i16a<u></u>reg64;<br>
+ break;<br>
+ case MVT::i32:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i32a<u></u>reg64;<br>
+ break;<br>
+ case MVT::i64:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i64a<u></u>reg64;<br>
+ break;<br>
+ case MVT::f32:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>f32a<u></u>reg64;<br>
+ break;<br>
+ case MVT::f64:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>f64a<u></u>reg64;<br>
+ break;<br>
+ }<br>
+ }<br>
+ break;<br>
case NVPTXISD::LDGV2:<br>
switch (EltVT.getSimpleVT().SimpleTy) {<br>
default:<br>
@@ -1374,6 +1618,55 @@ SDNode *NVPTXDAGToDAGISel::<u></u>SelectLDGL<u></u>DUV<br>
switch (N->getOpcode()) {<br>
default:<br>
return nullptr;<br>
+ case ISD::INTRINSIC_W_CHAIN:<br>
+ if (IsLDG) {<br>
+ switch (EltVT.getSimpleVT().SimpleTy) {<br>
+ default:<br>
+ return nullptr;<br>
+ case MVT::i8:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i8ar<u></u>eg;<br>
+ break;<br>
+ case MVT::i16:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i16a<u></u>reg;<br>
+ break;<br>
+ case MVT::i32:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i32a<u></u>reg;<br>
+ break;<br>
+ case MVT::i64:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i64a<u></u>reg;<br>
+ break;<br>
+ case MVT::f32:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>f32a<u></u>reg;<br>
+ break;<br>
+ case MVT::f64:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>f64a<u></u>reg;<br>
+ break;<br>
+ }<br>
+ } else {<br>
+ switch (EltVT.getSimpleVT().SimpleTy) {<br>
+ default:<br>
+ return nullptr;<br>
+ case MVT::i8:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i8ar<u></u>eg;<br>
+ break;<br>
+ case MVT::i16:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i16a<u></u>reg;<br>
+ break;<br>
+ case MVT::i32:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i32a<u></u>reg;<br>
+ break;<br>
+ case MVT::i64:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i64a<u></u>reg;<br>
+ break;<br>
+ case MVT::f32:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>f32a<u></u>reg;<br>
+ break;<br>
+ case MVT::f64:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>f64a<u></u>reg;<br>
+ break;<br>
+ }<br>
+ }<br>
+ break;<br>
case NVPTXISD::LDGV2:<br>
switch (EltVT.getSimpleVT().SimpleTy) {<br>
default:<br>
@@ -1466,7 +1759,7 @@ SDNode *NVPTXDAGToDAGISel::<u></u>SelectLDGL<u></u>DUV<br>
}<br>
<br>
MachineSDNode::mmo_iterator MemRefs0 = MF->allocateMemRefsArray(1);<br>
- MemRefs0[0] = cast<MemSDNode>(N)-><u></u>getMemOper<u></u>and();<br>
+ MemRefs0[0] = Mem->getMemOperand();<br>
cast<MachineSDNode>(LD)-><u></u>setM<u></u>emRefs(MemRefs0, MemRefs0 + 1);<br>
<br>
return LD;<br>
<br>
Modified: llvm/trunk/lib/Target/NVPTX/<u></u>NV<u></u>PTXISelDAGToDAG.h<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXISelDAGToDAG.h?rev=211939&r1=211938&r2=211939&view=diff" target="_blank">http://llvm.org/viewvc/llvm-<u></u>pr<u></u>oject/llvm/trunk/lib/Target/<u></u>NV<u></u>PTX/NVPTXISelDAGToDAG.h?rev=<u></u>21<u></u>1939&r1=211938&r2=211939&<u></u>view=<u></u>diff</a><br>
==============================<u></u><u></u>==============================<u></u><u></u>==================<br>
--- llvm/trunk/lib/Target/NVPTX/<u></u>NV<u></u>PTXISelDAGToDAG.h (original)<br>
+++ llvm/trunk/lib/Target/NVPTX/<u></u>NV<u></u>PTXISelDAGToDAG.h Fri Jun 27 13:35:51 2014<br>
@@ -59,10 +59,11 @@ private:<br>
<br>
SDNode *Select(SDNode *N) override;<br>
SDNode *SelectIntrinsicNoChain(SDNode *N);<br>
+ SDNode *SelectIntrinsicChain(SDNode *N);<br>
SDNode *SelectTexSurfHandle(SDNode *N);<br>
SDNode *SelectLoad(SDNode *N);<br>
SDNode *SelectLoadVector(SDNode *N);<br>
- SDNode *SelectLDGLDUVector(SDNode *N);<br>
+ SDNode *SelectLDGLDU(SDNode *N);<br>
SDNode *SelectStore(SDNode *N);<br>
SDNode *SelectStoreVector(SDNode *N);<br>
SDNode *SelectLoadParam(SDNode *N);<br>
<br>
Modified: llvm/trunk/lib/Target/NVPTX/<u></u>NV<u></u>PTXISelLowering.cpp<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXISelLowering.cpp?rev=211939&r1=211938&r2=211939&view=diff" target="_blank">http://llvm.org/viewvc/llvm-<u></u>pr<u></u>oject/llvm/trunk/lib/Target/<u></u>NV<u></u>PTX/NVPTXISelLowering.cpp?<u></u>rev=<u></u>211939&r1=211938&r2=<u></u>211939&<u></u>view=diff</a><br>
==============================<u></u><u></u>==============================<u></u><u></u>==================<br>
--- llvm/trunk/lib/Target/NVPTX/<u></u>NV<u></u>PTXISelLowering.cpp (original)<br>
+++ llvm/trunk/lib/Target/NVPTX/<u></u>NV<u></u>PTXISelLowering.cpp Fri Jun 27 13:35:51 2014<br>
@@ -2363,22 +2363,62 @@ bool NVPTXTargetLowering::<u></u>getTgtMem<u></u>Intri<br>
<br>
case Intrinsic::nvvm_ldu_global_i:<br>
case Intrinsic::nvvm_ldu_global_f:<br>
- case Intrinsic::nvvm_ldu_global_p:<br>
+ case Intrinsic::nvvm_ldu_global_p: {<br>
<br>
Info.opc = ISD::INTRINSIC_W_CHAIN;<br>
if (Intrinsic == Intrinsic::nvvm_ldu_global_i)<br>
Info.memVT = getValueType(I.getType());<br>
- else if (Intrinsic == Intrinsic::nvvm_ldu_global_p)<br>
+ else if(Intrinsic == Intrinsic::nvvm_ldu_global_p)<br>
+ Info.memVT = getPointerTy();<br>
+ else<br>
Info.memVT = getValueType(I.getType());<br>
+ Info.ptrVal = I.getArgOperand(0);<br>
+ Info.offset = 0;<br>
+ Info.vol = 0;<br>
+ Info.readMem = true;<br>
+ Info.writeMem = false;<br>
+<br>
+ // alignment is available as metadata.<br>
+ // Grab it and set the alignment.<br>
+ assert(I.<u></u>hasMetadataOtherThanD<u></u>ebugLoc() && "Must have alignment metadata");<br>
+ MDNode *AlignMD = I.getMetadata("align");<br>
+ assert(AlignMD && "Must have a non-null MDNode");<br>
+ assert(AlignMD-><u></u>getNumOperands<u></u>() == 1 && "Must have a single operand");<br>
+ Value *Align = AlignMD->getOperand(0);<br>
+ int64_t Alignment = cast<ConstantInt>(Align)-><u></u>getZ<u></u>ExtValue();<br>
+ Info.align = Alignment;<br>
+<br>
+ return true;<br>
+ }<br>
+ case Intrinsic::nvvm_ldg_global_i:<br>
+ case Intrinsic::nvvm_ldg_global_f:<br>
+ case Intrinsic::nvvm_ldg_global_p: {<br>
+<br>
+ Info.opc = ISD::INTRINSIC_W_CHAIN;<br>
+ if (Intrinsic == Intrinsic::nvvm_ldg_global_i)<br>
+ Info.memVT = getValueType(I.getType());<br>
+ else if(Intrinsic == Intrinsic::nvvm_ldg_global_p)<br>
+ Info.memVT = getPointerTy();<br>
else<br>
- Info.memVT = MVT::f32;<br>
+ Info.memVT = getValueType(I.getType());<br>
Info.ptrVal = I.getArgOperand(0);<br>
Info.offset = 0;<br>
Info.vol = 0;<br>
Info.readMem = true;<br>
Info.writeMem = false;<br>
- Info.align = 0;<br>
+<br>
+ // alignment is available as metadata.<br>
+ // Grab it and set the alignment.<br>
+ assert(I.<u></u>hasMetadataOtherThanD<u></u>ebugLoc() && "Must have alignment metadata");<br>
+ MDNode *AlignMD = I.getMetadata("align");<br>
+ assert(AlignMD && "Must have a non-null MDNode");<br>
+ assert(AlignMD-><u></u>getNumOperands<u></u>() == 1 && "Must have a single operand");<br>
+ Value *Align = AlignMD->getOperand(0);<br>
+ int64_t Alignment = cast<ConstantInt>(Align)-><u></u>getZ<u></u>ExtValue();<br>
+ Info.align = Alignment;<br>
+<br>
return true;<br>
+ }<br>
<br>
case Intrinsic::nvvm_tex_1d_v4f32_<u></u>i<u></u>32:<br>
case Intrinsic::nvvm_tex_1d_v4f32_<u></u>f<u></u>32:<br>
<br>
Modified: llvm/trunk/lib/Target/NVPTX/<u></u>NV<u></u>PTXIntrinsics.td<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXIntrinsics.td?rev=211939&r1=211938&r2=211939&view=diff" target="_blank">http://llvm.org/viewvc/llvm-<u></u>pr<u></u>oject/llvm/trunk/lib/Target/<u></u>NV<u></u>PTX/NVPTXIntrinsics.td?rev=<u></u>211<u></u>939&r1=211938&r2=211939&<u></u>view=<u></u>diff</a><br>
==============================<u></u><u></u>==============================<u></u><u></u>==================<br>
--- llvm/trunk/lib/Target/NVPTX/<u></u>NV<u></u>PTXIntrinsics.td (original)<br>
+++ llvm/trunk/lib/Target/NVPTX/<u></u>NV<u></u>PTXIntrinsics.td Fri Jun 27 13:35:51 2014<br>
@@ -1374,67 +1374,33 @@ def INT_PTX_SREG_WARPSIZE : F_SREG<"mov.<br>
// Support for ldu on sm_20 or later<br>
//---------------------------<u></u>-<u></u>-------<br>
<br>
-def ldu_i8 : PatFrag<(ops node:$ptr), (int_nvvm_ldu_global_i node:$ptr), [{<br>
- MemIntrinsicSDNode *M = cast<MemIntrinsicSDNode>(N);<br>
- return M->getMemoryVT() == MVT::i8;<br>
-}]>;<br>
-<br>
// Scalar<br>
-// @TODO: Revisit this, Changed imemAny to imem<br>
-multiclass LDU_G<string TyStr, NVPTXRegClass regclass, Intrinsic IntOp> {<br>
- def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src),<br>
- !strconcat("ldu.global.", TyStr),<br>
- [(set regclass:$result, (IntOp Int32Regs:$src))]>, Requires<[hasLDU]>;<br>
- def areg64: NVPTXInst<(outs regclass:$result), (ins Int64Regs:$src),<br>
- !strconcat("ldu.global.", TyStr),<br>
- [(set regclass:$result, (IntOp Int64Regs:$src))]>, Requires<[hasLDU]>;<br>
- def avar: NVPTXInst<(outs regclass:$result), (ins imem:$src),<br>
- !strconcat("ldu.global.", TyStr),<br>
- [(set regclass:$result, (IntOp (Wrapper tglobaladdr:$src)))]>,<br>
- Requires<[hasLDU]>;<br>
- def ari : NVPTXInst<(outs regclass:$result), (ins MEMri:$src),<br>
- !strconcat("ldu.global.", TyStr),<br>
- [(set regclass:$result, (IntOp ADDRri:$src))]>, Requires<[hasLDU]>;<br>
- def ari64 : NVPTXInst<(outs regclass:$result), (ins MEMri64:$src),<br>
- !strconcat("ldu.global.", TyStr),<br>
- [(set regclass:$result, (IntOp ADDRri64:$src))]>, Requires<[hasLDU]>;<br>
-}<br>
-<br>
-multiclass LDU_G_NOINTRIN<string TyStr, NVPTXRegClass regclass, PatFrag IntOp> {<br>
+multiclass LDU_G<string TyStr, NVPTXRegClass regclass> {<br>
def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src),<br>
!strconcat("ldu.global.", TyStr),<br>
- [(set regclass:$result, (IntOp Int32Regs:$src))]>, Requires<[hasLDU]>;<br>
+ []>, Requires<[hasLDU]>;<br>
def areg64: NVPTXInst<(outs regclass:$result), (ins Int64Regs:$src),<br>
!strconcat("ldu.global.", TyStr),<br>
- [(set regclass:$result, (IntOp Int64Regs:$src))]>, Requires<[hasLDU]>;<br>
- def avar: NVPTXInst<(outs regclass:$result), (ins imem:$src),<br>
+ []>, Requires<[hasLDU]>;<br>
+ def avar: NVPTXInst<(outs regclass:$result), (ins imemAny:$src),<br>
!strconcat("ldu.global.", TyStr),<br>
- [(set regclass:$result, (IntOp (Wrapper tglobaladdr:$src)))]>,<br>
- Requires<[hasLDU]>;<br>
+ []>, Requires<[hasLDU]>;<br>
def ari : NVPTXInst<(outs regclass:$result), (ins MEMri:$src),<br>
!strconcat("ldu.global.", TyStr),<br>
- [(set regclass:$result, (IntOp ADDRri:$src))]>, Requires<[hasLDU]>;<br>
+ []>, Requires<[hasLDU]>;<br>
def ari64 : NVPTXInst<(outs regclass:$result), (ins MEMri64:$src),<br>
!strconcat("ldu.global.", TyStr),<br>
- [(set regclass:$result, (IntOp ADDRri64:$src))]>, Requires<[hasLDU]>;<br>
+ []>, Requires<[hasLDU]>;<br>
}<br>
<br>
-defm INT_PTX_LDU_GLOBAL_i8 : LDU_G_NOINTRIN<"u8 \t$result, [$src];", Int16Regs,<br>
- ldu_i8>;<br>
-defm INT_PTX_LDU_GLOBAL_i16 : LDU_G<"u16 \t$result, [$src];", Int16Regs,<br>
-int_nvvm_ldu_global_i>;<br>
-defm INT_PTX_LDU_GLOBAL_i32 : LDU_G<"u32 \t$result, [$src];", Int32Regs,<br>
-int_nvvm_ldu_global_i>;<br>
-defm INT_PTX_LDU_GLOBAL_i64 : LDU_G<"u64 \t$result, [$src];", Int64Regs,<br>
-int_nvvm_ldu_global_i>;<br>
-defm INT_PTX_LDU_GLOBAL_f32 : LDU_G<"f32 \t$result, [$src];", Float32Regs,<br>
-int_nvvm_ldu_global_f>;<br>
-defm INT_PTX_LDU_GLOBAL_f64 : LDU_G<"f64 \t$result, [$src];", Float64Regs,<br>
-int_nvvm_ldu_global_f>;<br>
-defm INT_PTX_LDU_GLOBAL_p32 : LDU_G<"u32 \t$result, [$src];", Int32Regs,<br>
-int_nvvm_ldu_global_p>;<br>
-defm INT_PTX_LDU_GLOBAL_p64 : LDU_G<"u64 \t$result, [$src];", Int64Regs,<br>
-int_nvvm_ldu_global_p>;<br>
+defm INT_PTX_LDU_GLOBAL_i8 : LDU_G<"u8 \t$result, [$src];", Int16Regs>;<br>
+defm INT_PTX_LDU_GLOBAL_i16 : LDU_G<"u16 \t$result, [$src];", Int16Regs>;<br>
+defm INT_PTX_LDU_GLOBAL_i32 : LDU_G<"u32 \t$result, [$src];", Int32Regs>;<br>
+defm INT_PTX_LDU_GLOBAL_i64 : LDU_G<"u64 \t$result, [$src];", Int64Regs>;<br>
+defm INT_PTX_LDU_GLOBAL_f32 : LDU_G<"f32 \t$result, [$src];", Float32Regs>;<br>
+defm INT_PTX_LDU_GLOBAL_f64 : LDU_G<"f64 \t$result, [$src];", Float64Regs>;<br>
+defm INT_PTX_LDU_GLOBAL_p32 : LDU_G<"u32 \t$result, [$src];", Int32Regs>;<br>
+defm INT_PTX_LDU_GLOBAL_p64 : LDU_G<"u64 \t$result, [$src];", Int64Regs>;<br>
<br>
// vector<br>
<br>
@@ -1504,65 +1470,40 @@ defm INT_PTX_LDU_G_v4f32_ELE<br>
// Support for ldg on sm_35 or later<br>
//---------------------------<u></u>-<u></u>-------<br>
<br>
-def ldg_i8 : PatFrag<(ops node:$ptr), (int_nvvm_ldg_global_i node:$ptr), [{<br>
- MemIntrinsicSDNode *M = cast<MemIntrinsicSDNode>(N);<br>
- return M->getMemoryVT() == MVT::i8;<br>
-}]>;<br>
-<br>
-multiclass LDG_G<string TyStr, NVPTXRegClass regclass, Intrinsic IntOp> {<br>
- def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src),<br>
- !strconcat("<a href="http://ld.global.nc" target="_blank">ld.global.nc</a>.", TyStr),<br>
- [(set regclass:$result, (IntOp Int32Regs:$src))]>, Requires<[hasLDG]>;<br>
- def areg64: NVPTXInst<(outs regclass:$result), (ins Int64Regs:$src),<br>
- !strconcat("<a href="http://ld.global.nc" target="_blank">ld.global.nc</a>.", TyStr),<br>
- [(set regclass:$result, (IntOp Int64Regs:$src))]>, Requires<[hasLDG]>;<br>
- def avar: NVPTXInst<(outs regclass:$result), (ins imem:$src),<br>
- !strconcat("<a href="http://ld.global.nc" target="_blank">ld.global.nc</a>.", TyStr),<br>
- [(set regclass:$result, (IntOp (Wrapper tglobaladdr:$src)))]>,<br>
- Requires<[hasLDG]>;<br>
- def ari : NVPTXInst<(outs regclass:$result), (ins MEMri:$src),<br>
- !strconcat("<a href="http://ld.global.nc" target="_blank">ld.global.nc</a>.", TyStr),<br>
- [(set regclass:$result, (IntOp ADDRri:$src))]>, Requires<[hasLDG]>;<br>
- def ari64 : NVPTXInst<(outs regclass:$result), (ins MEMri64:$src),<br>
- !strconcat("<a href="http://ld.global.nc" target="_blank">ld.global.nc</a>.", TyStr),<br>
- [(set regclass:$result, (IntOp ADDRri64:$src))]>, Requires<[hasLDG]>;<br>
-}<br>
-<br>
-multiclass LDG_G_NOINTRIN<string TyStr, NVPTXRegClass regclass, PatFrag IntOp> {<br>
+multiclass LDG_G<string TyStr, NVPTXRegClass regclass> {<br>
def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src),<br>
!strconcat("<a href="http://ld.global.nc" target="_blank">ld.global.nc</a>.", TyStr),<br>
- [(set regclass:$result, (IntOp Int32Regs:$src))]>, Requires<[hasLDG]>;<br>
+ []>, Requires<[hasLDG]>;<br>
def areg64: NVPTXInst<(outs regclass:$result), (ins Int64Regs:$src),<br>
!strconcat("<a href="http://ld.global.nc" target="_blank">ld.global.nc</a>.", TyStr),<br>
- [(set regclass:$result, (IntOp Int64Regs:$src))]>, Requires<[hasLDG]>;<br>
- def avar: NVPTXInst<(outs regclass:$result), (ins imem:$src),<br>
+ []>, Requires<[hasLDG]>;<br>
+ def avar: NVPTXInst<(outs regclass:$result), (ins imemAny:$src),<br>
!strconcat("<a href="http://ld.global.nc" target="_blank">ld.global.nc</a>.", TyStr),<br>
- [(set regclass:$result, (IntOp (Wrapper tglobaladdr:$src)))]>,<br>
- Requires<[hasLDG]>;<br>
+ []>, Requires<[hasLDG]>;<br>
def ari : NVPTXInst<(outs regclass:$result), (ins MEMri:$src),<br>
!strconcat("<a href="http://ld.global.nc" target="_blank">ld.global.nc</a>.", TyStr),<br>
- [(set regclass:$result, (IntOp ADDRri:$src))]>, Requires<[hasLDG]>;<br>
+ []>, Requires<[hasLDG]>;<br>
def ari64 : NVPTXInst<(outs regclass:$result), (ins MEMri64:$src),<br>
!strconcat("<a href="http://ld.global.nc" target="_blank">ld.global.nc</a>.", TyStr),<br>
- [(set regclass:$result, (IntOp ADDRri64:$src))]>, Requires<[hasLDG]>;<br>
+ []>, Requires<[hasLDG]>;<br>
}<br>
<br>
defm INT_PTX_LDG_GLOBAL_i8<br>
- : LDG_G_NOINTRIN<"u8 \t$result, [$src];", Int16Regs, ldg_i8>;<br>
+ : LDG_G<"u8 \t$result, [$src];", Int16Regs>;<br>
defm INT_PTX_LDG_GLOBAL_i16<br>
- : LDG_G<"u16 \t$result, [$src];", Int16Regs, int_nvvm_ldg_global_i>;<br>
+ : LDG_G<"u16 \t$result, [$src];", Int16Regs>;<br>
defm INT_PTX_LDG_GLOBAL_i32<br>
- : LDG_G<"u32 \t$result, [$src];", Int32Regs, int_nvvm_ldg_global_i>;<br>
+ : LDG_G<"u32 \t$result, [$src];", Int32Regs>;<br>
defm INT_PTX_LDG_GLOBAL_i64<br>
- : LDG_G<"u64 \t$result, [$src];", Int64Regs, int_nvvm_ldg_global_i>;<br>
+ : LDG_G<"u64 \t$result, [$src];", Int64Regs>;<br>
defm INT_PTX_LDG_GLOBAL_f32<br>
- : LDG_G<"f32 \t$result, [$src];", Float32Regs, int_nvvm_ldg_global_f>;<br>
+ : LDG_G<"f32 \t$result, [$src];", Float32Regs>;<br>
defm INT_PTX_LDG_GLOBAL_f64<br>
- : LDG_G<"f64 \t$result, [$src];", Float64Regs, int_nvvm_ldg_global_f>;<br>
+ : LDG_G<"f64 \t$result, [$src];", Float64Regs>;<br>
defm INT_PTX_LDG_GLOBAL_p32<br>
- : LDG_G<"u32 \t$result, [$src];", Int32Regs, int_nvvm_ldg_global_p>;<br>
+ : LDG_G<"u32 \t$result, [$src];", Int32Regs>;<br>
defm INT_PTX_LDG_GLOBAL_p64<br>
- : LDG_G<"u64 \t$result, [$src];", Int64Regs, int_nvvm_ldg_global_p>;<br>
+ : LDG_G<"u64 \t$result, [$src];", Int64Regs>;<br>
<br>
// vector<br>
<br>
<br>
Modified: llvm/trunk/test/CodeGen/NVPTX/<u></u><u></u>ldu-i8.ll<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/NVPTX/ldu-i8.ll?rev=211939&r1=211938&r2=211939&view=diff" target="_blank">http://llvm.org/viewvc/llvm-<u></u>pr<u></u>oject/llvm/trunk/test/<u></u>CodeGen/<u></u>NVPTX/ldu-i8.ll?rev=<u></u>211939&r1=<u></u>211938&r2=211939&<u></u>view=diff</a><br>
==============================<u></u><u></u>==============================<u></u><u></u>==================<br>
--- llvm/trunk/test/CodeGen/NVPTX/<u></u><u></u>ldu-i8.ll (original)<br>
+++ llvm/trunk/test/CodeGen/NVPTX/<u></u><u></u>ldu-i8.ll Fri Jun 27 13:35:51 2014<br>
@@ -2,13 +2,15 @@<br>
<br>
target datalayout = "e-p:32:32:32-i1:8:8-i8:8:8-<u></u>i1<u></u>6:16:16-i32:32:32-i64:64:64-<u></u>f3<u></u>2:32:32-f64:64:64-v16:16:16-<u></u>v3<u></u>2:32:32-v64:64:64-v128:128:<u></u>128<u></u>-n16:32:64"<br>
<br>
-declare i8 @llvm.nvvm.ldu.global.i.i8(i8*<u></u><u></u>)<br>
+declare i8 @llvm.nvvm.ldu.global.i.i8.<u></u>p0i<u></u>8(i8*)<br>
<br>
define i8 @foo(i8* %a) {<br>
; Ensure we properly truncate off the high-order 24 bits<br>
; CHECK: ldu.global.u8<br>
; CHECK: cvt.u32.u16<br>
; CHECK: and.b32 %r{{[0-9]+}}, %r{{[0-9]+}}, 255<br>
- %val = tail call i8 @llvm.nvvm.ldu.global.i.i8(i8* %a)<br>
+ %val = tail call i8 @llvm.nvvm.ldu.global.i.i8.<u></u>p0i<u></u>8(i8* %a), !align !0<br>
ret i8 %val<br>
}<br>
+<br>
+!0 = metadata !{i32 4}<br>
<br>
Added: llvm/trunk/test/CodeGen/NVPTX/<u></u><u></u>ldu-ldg.ll<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/NVPTX/ldu-ldg.ll?rev=211939&view=auto" target="_blank">http://llvm.org/viewvc/llvm-<u></u>pr<u></u>oject/llvm/trunk/test/<u></u>CodeGen/<u></u>NVPTX/ldu-ldg.ll?rev=<u></u>211939&<u></u>view=auto</a><br>
==============================<u></u><u></u>==============================<u></u><u></u>==================<br>
--- llvm/trunk/test/CodeGen/NVPTX/<u></u><u></u>ldu-ldg.ll (added)<br>
+++ llvm/trunk/test/CodeGen/NVPTX/<u></u><u></u>ldu-ldg.ll Fri Jun 27 13:35:51 2014<br>
@@ -0,0 +1,40 @@<br>
+; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s<br>
+<br>
+<br>
+declare i8 @llvm.nvvm.ldu.global.i.i8.<u></u>p1i<u></u>8(i8 addrspace(1)* %ptr)<br>
+declare i32 @llvm.nvvm.ldu.global.i.i32.<u></u>p1<u></u>i32(i32 addrspace(1)* %ptr)<br>
+declare i8 @llvm.nvvm.ldg.global.i.i8.<u></u>p1i<u></u>8(i8 addrspace(1)* %ptr)<br>
+declare i32 @llvm.nvvm.ldg.global.i.i32.<u></u>p1<u></u>i32(i32 addrspace(1)* %ptr)<br>
+<br>
+<br>
+; CHECK: func0<br>
+define i8 @func0(i8 addrspace(1)* %ptr) {<br>
+; ldu.global.u8<br>
+ %val = tail call i8 @llvm.nvvm.ldu.global.i.i8.<u></u>p1i<u></u>8(i8 addrspace(1)* %ptr), !align !0<br>
+ ret i8 %val<br>
+}<br>
+<br>
+; CHECK: func1<br>
+define i32 @func1(i32 addrspace(1)* %ptr) {<br>
+; ldu.global.u32<br>
+ %val = tail call i32 @llvm.nvvm.ldu.global.i.i32.<u></u>p1<u></u>i32(i32 addrspace(1)* %ptr), !align !0<br>
+ ret i32 %val<br>
+}<br>
+<br>
+; CHECK: func2<br>
+define i8 @func2(i8 addrspace(1)* %ptr) {<br>
+; ld.global.nc.u8<br>
+ %val = tail call i8 @llvm.nvvm.ldg.global.i.i8.<u></u>p1i<u></u>8(i8 addrspace(1)* %ptr), !align !0<br>
+ ret i8 %val<br>
+}<br>
+<br>
+; CHECK: func3<br>
+define i32 @func3(i32 addrspace(1)* %ptr) {<br>
+; ld.global.nc.u32<br>
+ %val = tail call i32 @llvm.nvvm.ldg.global.i.i32.<u></u>p1<u></u>i32(i32 addrspace(1)* %ptr), !align !0<br>
+ ret i32 %val<br>
+}<br>
+<br>
+<br>
+<br>
+!0 = metadata !{i32 4}<br>
<br>
Modified: llvm/trunk/test/CodeGen/NVPTX/<u></u><u></u>ldu-reg-plus-offset.ll<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/NVPTX/ldu-reg-plus-offset.ll?rev=211939&r1=211938&r2=211939&view=diff" target="_blank">http://llvm.org/viewvc/llvm-<u></u>pr<u></u>oject/llvm/trunk/test/<u></u>CodeGen/<u></u>NVPTX/ldu-reg-plus-<u></u>offset.ll?<u></u>rev=211939&r1=<u></u>211938&r2=<u></u>211939&view=diff</a><br>
==============================<u></u><u></u>==============================<u></u><u></u>==================<br>
--- llvm/trunk/test/CodeGen/NVPTX/<u></u><u></u>ldu-reg-plus-offset.ll (original)<br>
+++ llvm/trunk/test/CodeGen/NVPTX/<u></u><u></u>ldu-reg-plus-offset.ll Fri Jun 27 13:35:51 2014<br>
@@ -7,9 +7,9 @@ define void @reg_plus_offset(i32* %a) {<br>
; CHECK: ldu.global.u32 %r{{[0-9]+}}, [%r{{[0-9]+}}+32];<br>
; CHECK: ldu.global.u32 %r{{[0-9]+}}, [%r{{[0-9]+}}+36];<br>
%p2 = getelementptr i32* %a, i32 8<br>
- %t1 = call i32 @llvm.nvvm.ldu.global.i.i32(<u></u>i3<u></u>2* %p2), !align !1<br>
+ %t1 = call i32 @llvm.nvvm.ldu.global.i.i32.<u></u>p0<u></u>i32(i32* %p2), !align !1<br>
%p3 = getelementptr i32* %a, i32 9<br>
- %t2 = call i32 @llvm.nvvm.ldu.global.i.i32(<u></u>i3<u></u>2* %p3), !align !1<br>
+ %t2 = call i32 @llvm.nvvm.ldu.global.i.i32.<u></u>p0<u></u>i32(i32* %p3), !align !1<br>
%t3 = mul i32 %t1, %t2<br>
store i32 %t3, i32* %a<br>
ret void<br>
@@ -17,5 +17,5 @@ define void @reg_plus_offset(i32* %a) {<br>
<br>
!1 = metadata !{ i32 4 }<br>
<br>
-declare i32 @llvm.nvvm.ldu.global.i.i32(<u></u>i3<u></u>2*)<br>
+declare i32 @llvm.nvvm.ldu.global.i.i32.<u></u>p0<u></u>i32(i32*)<br>
declare i32 @llvm.nvvm.read.ptx.sreg.tid.<u></u>x<u></u>()<br>
<br>
<br>
______________________________<u></u><u></u>_________________<br>
llvm-commits mailing list<br>
<a href="mailto:llvm-commits@cs.uiuc.edu" target="_blank">llvm-commits@cs.uiuc.edu</a><br>
<a href="http://lists.cs.uiuc.edu/mailman/listinfo/llvm-commits" target="_blank">http://lists.cs.uiuc.edu/<u></u>mailm<u></u>an/listinfo/llvm-commits</a><br>
</blockquote></div>