Hi Justin, <br><br>I sent out a patch (<a href="http://reviews.llvm.org/D5093">http://reviews.llvm.org/D5093</a>) that makes the alignment an argument to ldu/ldg intrinsics. PTAL. <div><br></div><div>Thanks,</div><div>Jingyue</div>
<br><div class="gmail_quote">On Mon Aug 25 2014 at 9:30:57 AM Jingyue Wu <<a href="mailto:jingyue@google.com">jingyue@google.com</a>> wrote:<br><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
I agree. That would cut off the issue of metadata being removed by optimizations. <div><br></div><div>Does the current NVPTX backend use or check the alignment at all? I tried, for example, using llvm.nvvm.ldg.global.f.v4f32.<u></u>p0v4f32 with alignment 4, and the emitted PTX still uses ld.global.nc.v4.f32 even though the alignment is not 4 * sizeof(float). </div>
<div><br></div><div>Jingyue</div><div><br></div><div><div class="gmail_quote">On Fri Aug 22 2014 at 8:54:47 AM 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">
<div>
<div style="font-size:10pt;color:#000000;background-color:#ffffff;font-family:'Lucida Console',Monaco,monospace">
<p>Hi Jingyue,</p>
<p><br>
</p>
<p>I agree that this is fragile and needs to be reworked. "Officially", the ldu/ldg intrinsics are not meant to be used by front-ends (they do not exist in the NVVM IR spec). In libnvvm, they are generated automatically through IR transformations just before
PTX code generation (and after all of the 'opt'-level transformations). However, for upstream NVPTX, it makes a lot of sense to expose these intrinsics so users can generate them, especially since the ldg/ldu transforamtion passes are not upstream. I believe
the proper way to handle these would be to make the alignment an explicit argument, similar to llvm.memset/llvm.memcpy/llvm.<u></u>m<u></u>emmove. Do you agree?<br>
</p>
<p><br>
</p>
<p><br>
</p>
<div style="color:#282828">
<hr style="display:inline-block;width:98%">
<div dir="ltr"><font face="Calibri, sans-serif" color="#000000" style="font-size:11pt"><b>From:</b> Jingyue Wu <<a href="mailto:jingyue@google.com" target="_blank">jingyue@google.com</a>><br>
<b>Sent:</b> Thursday, August 21, 2014 6:59 PM<br>
<b>To:</b> Justin Holewinski<br>
<b>Cc:</b> <a href="mailto:llvm-commits@cs.uiuc.edu" target="_blank">llvm-commits@cs.uiuc.edu</a>; Eli Bendersky; Mark Heffernan; Cong Hou<br>
<b>Subject:</b> Re: [llvm] r211939 - [NVPTX] Fix handling of ldg/ldu intrinsics.</font>
<div> </div>
</div></div></div></div><div><div style="font-size:10pt;color:#000000;background-color:#ffffff;font-family:'Lucida Console',Monaco,monospace"><div style="color:#282828">
<div>Hi Justin, <br>
<br>
The assertion on NVPTXISelLowering.cpp:3305
<div><br>
</div>
<div>assert(I.<u></u>hasMetadataOtherThanD<u></u>ebugLoc() && "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:0px 0px 0px 0.8ex;border-left-width:1px;border-left-color:#cccccc;border-left-style: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><u></u><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><u></u><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><u></u><u></u>ldu-ldg.ll<br>
Modified:<br>
llvm/trunk/include/llvm/IR/<u></u>Int<u></u><u></u><u></u>rinsicsNVVM.td<br>
llvm/trunk/lib/Target/NVPTX/<u></u>NV<u></u><u></u><u></u>PTXISelDAGToDAG.cpp<br>
llvm/trunk/lib/Target/NVPTX/<u></u>NV<u></u><u></u><u></u>PTXISelDAGToDAG.h<br>
llvm/trunk/lib/Target/NVPTX/<u></u>NV<u></u><u></u><u></u>PTXISelLowering.cpp<br>
llvm/trunk/lib/Target/NVPTX/<u></u>NV<u></u><u></u><u></u>PTXIntrinsics.td<br>
llvm/trunk/test/CodeGen/NVPTX/<u></u><u></u><u></u><u></u>ldu-i8.ll<br>
llvm/trunk/test/CodeGen/NVPTX/<u></u><u></u><u></u><u></u>ldu-reg-plus-offset.ll<br>
<br>
Modified: llvm/trunk/include/llvm/IR/<u></u>Int<u></u><u></u><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><u></u><u></u>oject/llvm/trunk/include/<u></u>llvm/<u></u><u></u><u></u>IR/IntrinsicsNVVM.td?rev=<u></u>21193<u></u><u></u><u></u>9&r1=211938&r2=211939&<u></u>view=<u></u>dif<u></u><u></u>f</a><br>
==============================<u></u><u></u><u></u><u></u>==============================<u></u><u></u><u></u><u></u>==================<br>
--- llvm/trunk/include/llvm/IR/<u></u>Int<u></u><u></u><u></u>rinsicsNVVM.td (original)<br>
+++ llvm/trunk/include/llvm/IR/<u></u>Int<u></u><u></u><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><u></u><u></u><0>>], [IntrReadMem, NoCapture<0>],<br>
+ [LLVMAnyPointerType<<u></u>LLVMMatchT<u></u><u></u><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><u></u><u></u><0>>], [IntrReadMem, NoCapture<0>],<br>
+ [LLVMAnyPointerType<<u></u>LLVMMatchT<u></u><u></u><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><u></u><u></u><0>>], [IntrReadMem, NoCapture<0>],<br>
+ [LLVMAnyPointerType<<u></u>LLVMMatchT<u></u><u></u><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><u></u><u></u><0>>], [IntrReadMem, NoCapture<0>],<br>
+ [LLVMAnyPointerType<<u></u>LLVMMatchT<u></u><u></u><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><u></u><u></u><0>>], [IntrReadMem, NoCapture<0>],<br>
+ [LLVMAnyPointerType<<u></u>LLVMMatchT<u></u><u></u><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><u></u><u></u><0>>], [IntrReadMem, NoCapture<0>],<br>
+ [LLVMAnyPointerType<<u></u>LLVMMatchT<u></u><u></u><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><u></u><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><u></u><u></u>oject/llvm/trunk/lib/Target/<u></u>NV<u></u><u></u><u></u>PTX/NVPTXISelDAGToDAG.cpp?<u></u>rev=<u></u><u></u><u></u>211939&r1=211938&r2=<u></u>211939&<u></u>vie<u></u><u></u>w=diff</a><br>
==============================<u></u><u></u><u></u><u></u>==============================<u></u><u></u><u></u><u></u>==================<br>
--- llvm/trunk/lib/Target/NVPTX/<u></u>NV<u></u><u></u><u></u>PTXISelDAGToDAG.cpp (original)<br>
+++ llvm/trunk/lib/Target/NVPTX/<u></u>NV<u></u><u></u><u></u>PTXISelDAGToDAG.cpp Fri Jun 27 13:35:51 2014<br>
@@ -141,7 +141,7 @@ SDNode *NVPTXDAGToDAGISel::Select(<u></u>SDN<u></u><u></u><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><u></u><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><u></u><u></u>:<br>
@@ -273,6 +276,21 @@ SDNode *NVPTXDAGToDAGISel::Select(<u></u>SDN<u></u><u></u><u></u>ode<br>
return SelectCode(N);<br>
}<br>
<br>
+SDNode *NVPTXDAGToDAGISel::<u></u>SelectIntr<u></u><u></u><u></u>insicChain(SDNode *N) {<br>
+ unsigned IID = cast<ConstantSDNode>(N-><u></u>getOpe<u></u><u></u><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><u></u><u></u>;<br>
@@ -990,22 +1008,101 @@ SDNode *NVPTXDAGToDAGISel::<u></u>SelectLoad<u></u><u></u><u></u>Vec<br>
return LD;<br>
}<br>
<br>
-SDNode *NVPTXDAGToDAGISel::<u></u>SelectLDGL<u></u><u></u><u></u>DUVector(SDNode *N) {<br>
+SDNode *NVPTXDAGToDAGISel::<u></u>SelectLDGL<u></u><u></u><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><u></u><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><u></u><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><u></u><u></u>ar;<br>
+ break;<br>
+ case MVT::i16:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i16a<u></u><u></u><u></u>var;<br>
+ break;<br>
+ case MVT::i32:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i32a<u></u><u></u><u></u>var;<br>
+ break;<br>
+ case MVT::i64:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i64a<u></u><u></u><u></u>var;<br>
+ break;<br>
+ case MVT::f32:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>f32a<u></u><u></u><u></u>var;<br>
+ break;<br>
+ case MVT::f64:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>f64a<u></u><u></u><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><u></u><u></u>ar;<br>
+ break;<br>
+ case MVT::i16:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i16a<u></u><u></u><u></u>var;<br>
+ break;<br>
+ case MVT::i32:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i32a<u></u><u></u><u></u>var;<br>
+ break;<br>
+ case MVT::i64:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i64a<u></u><u></u><u></u>var;<br>
+ break;<br>
+ case MVT::f32:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>f32a<u></u><u></u><u></u>var;<br>
+ break;<br>
+ case MVT::f64:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>f64a<u></u><u></u><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><u></u><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><u></u><u></u>i64;<br>
+ break;<br>
+ case MVT::i16:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i16a<u></u><u></u><u></u>ri64;<br>
+ break;<br>
+ case MVT::i32:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i32a<u></u><u></u><u></u>ri64;<br>
+ break;<br>
+ case MVT::i64:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i64a<u></u><u></u><u></u>ri64;<br>
+ break;<br>
+ case MVT::f32:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>f32a<u></u><u></u><u></u>ri64;<br>
+ break;<br>
+ case MVT::f64:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>f64a<u></u><u></u><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><u></u><u></u>i64;<br>
+ break;<br>
+ case MVT::i16:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i16a<u></u><u></u><u></u>ri64;<br>
+ break;<br>
+ case MVT::i32:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i32a<u></u><u></u><u></u>ri64;<br>
+ break;<br>
+ case MVT::i64:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i64a<u></u><u></u><u></u>ri64;<br>
+ break;<br>
+ case MVT::f32:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>f32a<u></u><u></u><u></u>ri64;<br>
+ break;<br>
+ case MVT::f64:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>f64a<u></u><u></u><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><u></u><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><u></u><u></u>i;<br>
+ break;<br>
+ case MVT::i16:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i16a<u></u><u></u><u></u>ri;<br>
+ break;<br>
+ case MVT::i32:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i32a<u></u><u></u><u></u>ri;<br>
+ break;<br>
+ case MVT::i64:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i64a<u></u><u></u><u></u>ri;<br>
+ break;<br>
+ case MVT::f32:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>f32a<u></u><u></u><u></u>ri;<br>
+ break;<br>
+ case MVT::f64:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>f64a<u></u><u></u><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><u></u><u></u>i;<br>
+ break;<br>
+ case MVT::i16:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i16a<u></u><u></u><u></u>ri;<br>
+ break;<br>
+ case MVT::i32:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i32a<u></u><u></u><u></u>ri;<br>
+ break;<br>
+ case MVT::i64:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i64a<u></u><u></u><u></u>ri;<br>
+ break;<br>
+ case MVT::f32:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>f32a<u></u><u></u><u></u>ri;<br>
+ break;<br>
+ case MVT::f64:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>f64a<u></u><u></u><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><u></u><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><u></u><u></u>eg64;<br>
+ break;<br>
+ case MVT::i16:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i16a<u></u><u></u><u></u>reg64;<br>
+ break;<br>
+ case MVT::i32:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i32a<u></u><u></u><u></u>reg64;<br>
+ break;<br>
+ case MVT::i64:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i64a<u></u><u></u><u></u>reg64;<br>
+ break;<br>
+ case MVT::f32:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>f32a<u></u><u></u><u></u>reg64;<br>
+ break;<br>
+ case MVT::f64:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>f64a<u></u><u></u><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><u></u><u></u>eg64;<br>
+ break;<br>
+ case MVT::i16:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i16a<u></u><u></u><u></u>reg64;<br>
+ break;<br>
+ case MVT::i32:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i32a<u></u><u></u><u></u>reg64;<br>
+ break;<br>
+ case MVT::i64:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i64a<u></u><u></u><u></u>reg64;<br>
+ break;<br>
+ case MVT::f32:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>f32a<u></u><u></u><u></u>reg64;<br>
+ break;<br>
+ case MVT::f64:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>f64a<u></u><u></u><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><u></u><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><u></u><u></u>eg;<br>
+ break;<br>
+ case MVT::i16:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i16a<u></u><u></u><u></u>reg;<br>
+ break;<br>
+ case MVT::i32:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i32a<u></u><u></u><u></u>reg;<br>
+ break;<br>
+ case MVT::i64:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>i64a<u></u><u></u><u></u>reg;<br>
+ break;<br>
+ case MVT::f32:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>f32a<u></u><u></u><u></u>reg;<br>
+ break;<br>
+ case MVT::f64:<br>
+ Opcode = NVPTX::INT_PTX_LDG_GLOBAL_<u></u>f64a<u></u><u></u><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><u></u><u></u>eg;<br>
+ break;<br>
+ case MVT::i16:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i16a<u></u><u></u><u></u>reg;<br>
+ break;<br>
+ case MVT::i32:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i32a<u></u><u></u><u></u>reg;<br>
+ break;<br>
+ case MVT::i64:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>i64a<u></u><u></u><u></u>reg;<br>
+ break;<br>
+ case MVT::f32:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>f32a<u></u><u></u><u></u>reg;<br>
+ break;<br>
+ case MVT::f64:<br>
+ Opcode = NVPTX::INT_PTX_LDU_GLOBAL_<u></u>f64a<u></u><u></u><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><u></u><u></u>DUV<br>
}<br>
<br>
MachineSDNode::mmo_iterator MemRefs0 = MF->allocateMemRefsArray(1);<br>
- MemRefs0[0] = cast<MemSDNode>(N)-><u></u>getMemOper<u></u><u></u><u></u>and();<br>
+ MemRefs0[0] = Mem->getMemOperand();<br>
cast<MachineSDNode>(LD)-><u></u>setM<u></u><u></u><u></u>emRefs(MemRefs0, MemRefs0 + 1);<br>
<br>
return LD;<br>
<br>
Modified: llvm/trunk/lib/Target/NVPTX/<u></u>NV<u></u><u></u><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><u></u><u></u>oject/llvm/trunk/lib/Target/<u></u>NV<u></u><u></u><u></u>PTX/NVPTXISelDAGToDAG.h?rev=<u></u>21<u></u><u></u><u></u>1939&r1=211938&r2=211939&<u></u>view=<u></u><u></u><u></u>diff</a><br>
==============================<u></u><u></u><u></u><u></u>==============================<u></u><u></u><u></u><u></u>==================<br>
--- llvm/trunk/lib/Target/NVPTX/<u></u>NV<u></u><u></u><u></u>PTXISelDAGToDAG.h (original)<br>
+++ llvm/trunk/lib/Target/NVPTX/<u></u>NV<u></u><u></u><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><u></u><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><u></u><u></u>oject/llvm/trunk/lib/Target/<u></u>NV<u></u><u></u><u></u>PTX/NVPTXISelLowering.cpp?<u></u>rev=<u></u><u></u><u></u>211939&r1=211938&r2=<u></u>211939&<u></u>vie<u></u><u></u>w=diff</a><br>
==============================<u></u><u></u><u></u><u></u>==============================<u></u><u></u><u></u><u></u>==================<br>
--- llvm/trunk/lib/Target/NVPTX/<u></u>NV<u></u><u></u><u></u>PTXISelLowering.cpp (original)<br>
+++ llvm/trunk/lib/Target/NVPTX/<u></u>NV<u></u><u></u><u></u>PTXISelLowering.cpp Fri Jun 27 13:35:51 2014<br>
@@ -2363,22 +2363,62 @@ bool NVPTXTargetLowering::<u></u>getTgtMem<u></u><u></u><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><u></u><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><u></u><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><u></u><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><u></u><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><u></u><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><u></u><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><u></u><u></u>32:<br>
case Intrinsic::nvvm_tex_1d_v4f32_<u></u>f<u></u><u></u><u></u>32:<br>
<br>
Modified: llvm/trunk/lib/Target/NVPTX/<u></u>NV<u></u><u></u><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><u></u><u></u>oject/llvm/trunk/lib/Target/<u></u>NV<u></u><u></u><u></u>PTX/NVPTXIntrinsics.td?rev=<u></u>211<u></u><u></u><u></u>939&r1=211938&r2=211939&<u></u>view=<u></u>d<u></u><u></u>iff</a><br>
==============================<u></u><u></u><u></u><u></u>==============================<u></u><u></u><u></u><u></u>==================<br>
--- llvm/trunk/lib/Target/NVPTX/<u></u>NV<u></u><u></u><u></u>PTXIntrinsics.td (original)<br>
+++ llvm/trunk/lib/Target/NVPTX/<u></u>NV<u></u><u></u><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><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><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><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><u></u><u></u>oject/llvm/trunk/test/<u></u>CodeGen/<u></u><u></u><u></u>NVPTX/ldu-i8.ll?rev=<u></u>211939&r1=<u></u><u></u><u></u>211938&r2=211939&<u></u>view=diff</a><br>
==============================<u></u><u></u><u></u><u></u>==============================<u></u><u></u><u></u><u></u>==================<br>
--- llvm/trunk/test/CodeGen/NVPTX/<u></u><u></u><u></u><u></u>ldu-i8.ll (original)<br>
+++ llvm/trunk/test/CodeGen/NVPTX/<u></u><u></u><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><u></u><u></u>6:16:16-i32:32:32-i64:64:64-<u></u>f3<u></u><u></u><u></u>2:32:32-f64:64:64-v16:16:16-<u></u>v3<u></u><u></u><u></u>2:32:32-v64:64:64-v128:128:<u></u>128<u></u><u></u><u></u>-n16:32:64"<br>
<br>
-declare i8 @llvm.nvvm.ldu.global.i.i8(i8*<u></u><u></u><u></u><u></u>)<br>
+declare i8 @llvm.nvvm.ldu.global.i.i8.<u></u>p0i<u></u><u></u><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><u></u><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><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><u></u><u></u>oject/llvm/trunk/test/<u></u>CodeGen/<u></u><u></u><u></u>NVPTX/ldu-ldg.ll?rev=<u></u>211939&<u></u>vi<u></u><u></u>ew=auto</a><br>
==============================<u></u><u></u><u></u><u></u>==============================<u></u><u></u><u></u><u></u>==================<br>
--- llvm/trunk/test/CodeGen/NVPTX/<u></u><u></u><u></u><u></u>ldu-ldg.ll (added)<br>
+++ llvm/trunk/test/CodeGen/NVPTX/<u></u><u></u><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><u></u><u></u>8(i8 addrspace(1)* %ptr)<br>
+declare i32 @llvm.nvvm.ldu.global.i.i32.<u></u>p1<u></u><u></u><u></u>i32(i32 addrspace(1)* %ptr)<br>
+declare i8 @llvm.nvvm.ldg.global.i.i8.<u></u>p1i<u></u><u></u><u></u>8(i8 addrspace(1)* %ptr)<br>
+declare i32 @llvm.nvvm.ldg.global.i.i32.<u></u>p1<u></u><u></u><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><u></u><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><u></u><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><u></u><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><u></u><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><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><u></u><u></u>oject/llvm/trunk/test/<u></u>CodeGen/<u></u><u></u><u></u>NVPTX/ldu-reg-plus-<u></u>offset.ll?<u></u>r<u></u><u></u>ev=211939&r1=<u></u>211938&r2=<u></u>211939&<u></u><u></u>view=diff</a><br>
==============================<u></u><u></u><u></u><u></u>==============================<u></u><u></u><u></u><u></u>==================<br>
--- llvm/trunk/test/CodeGen/NVPTX/<u></u><u></u><u></u><u></u>ldu-reg-plus-offset.ll (original)<br>
+++ llvm/trunk/test/CodeGen/NVPTX/<u></u><u></u><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><u></u><u></u>2* %p2), !align !1<br>
+ %t1 = call i32 @llvm.nvvm.ldu.global.i.i32.<u></u>p0<u></u><u></u><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><u></u><u></u>2* %p3), !align !1<br>
+ %t2 = call i32 @llvm.nvvm.ldu.global.i.i32.<u></u>p0<u></u><u></u><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><u></u><u></u>2*)<br>
+declare i32 @llvm.nvvm.ldu.global.i.i32.<u></u>p0<u></u><u></u><u></u>i32(i32*)<br>
declare i32 @llvm.nvvm.read.ptx.sreg.tid.<u></u>x<u></u><u></u><u></u>()<br>
<br>
<br>
______________________________<u></u><u></u><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><u></u><u></u>an/listinfo/llvm-commits</a><br>
</blockquote>
</div>
</div>
</div></div></div><div><div style="font-size:10pt;color:#000000;background-color:#ffffff;font-family:'Lucida Console',Monaco,monospace"><div style="color:#282828"></div>
</div>
<div>
<hr>
</div>
<div>This email message is for the sole use of the intended recipient(s) and may
contain confidential information. Any unauthorized review, use, disclosure
or distribution is prohibited. If you are not the intended recipient,
please contact the sender by reply email and destroy all copies of the original
message. </div>
<div>
<hr>
</div>
</div>
</blockquote></div></div></blockquote></div>