<div class="gmail_quote">On Mon, Nov 14, 2011 at 12:55 PM, Villmow, Micah <span dir="ltr"><<a href="mailto:Micah.Villmow@amd.com">Micah.Villmow@amd.com</a>></span> wrote:<br><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex;">
<div lang="EN-US" link="blue" vlink="purple">
<div>
<p class="MsoNormal"><span style="font-size:11.0pt;color:#1F497D">Justin,<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;color:#1F497D">Add this to your TargetLowering constructor, this fixes the mem* issue.<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;color:#1F497D"><u></u> <u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;color:#1F497D"> maxStoresPerMemcpy = 4096;<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;color:#1F497D"> maxStoresPerMemmove = 4096;<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;color:#1F497D"> maxStoresPerMemset = 4096;</span></p></div></div></blockquote><div><br></div><div>Thanks for this! I applied it in r144551. However, this particular kernel still will not pass through on account of the lack of support to sign-extend loads from i8 to i32.</div>
<div> </div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex;"><div lang="EN-US" link="blue" vlink="purple"><div><p class="MsoNormal"><span style="font-size:11.0pt;color:#1F497D"><u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;color:#1F497D"><u></u> <u></u></span></p>
<div style="border:none;border-left:solid blue 1.5pt;padding:0in 0in 0in 4.0pt">
<div>
<div style="border:none;border-top:solid #B5C4DF 1.0pt;padding:3.0pt 0in 0in 0in">
<p class="MsoNormal"><b><span style="font-size:10.0pt">From:</span></b><span style="font-size:10.0pt"> <a href="mailto:llvmdev-bounces@cs.uiuc.edu" target="_blank">llvmdev-bounces@cs.uiuc.edu</a> [mailto:<a href="mailto:llvmdev-bounces@cs.uiuc.edu" target="_blank">llvmdev-bounces@cs.uiuc.edu</a>]
<b>On Behalf Of </b>Justin Holewinski<br>
<b>Sent:</b> Monday, November 14, 2011 7:12 AM<br>
<b>To:</b> Alberto Magni<br>
<b>Cc:</b> <a href="mailto:llvmdev@cs.uiuc.edu" target="_blank">llvmdev@cs.uiuc.edu</a><br>
<b>Subject:</b> Re: [LLVMdev] PTX backend fatal error<u></u><u></u></span></p>
</div>
</div><div><div class="adm"><div id="q_133a34b9b8e5f002_1" class="ajR h4"><div class="ajT"></div></div></div><div class="h5">
<p class="MsoNormal"><u></u> <u></u></p>
<div>
<p class="MsoNormal">On Mon, Nov 14, 2011 at 8:57 AM, Alberto Magni <<a href="mailto:alberto.magni86@gmail.com" target="_blank">alberto.magni86@gmail.com</a>> wrote:<u></u><u></u></p>
<p class="MsoNormal">Hi everybody,<br>
<br>
I am testing the PTX backend using the OpenCL NVIDIA SDK benchmarks.<br>
Compiling the Histogram64.cl program I get a several backend errors.<br>
<br>
I isolated one of them in the following kernel program:<br>
<br>
__kernel void kernel_function(__global int *input) {<br>
__local char localArray[16];<br>
for(unsigned int index = 0; index < 16; ++index)<br>
localArray[index] = 0;<br>
input[0] = localArray[get_local_id(0)];<br>
}<br>
<br>
fatal error: error in backend: Cannot select:<br>
0x5810cc0: i32,ch = load 0x57fa148,<br>
0x5810ac0, 0x58105c0<LD1[%arrayidx1], sext<br>
from i8> [ID=9]<br>
0x5810ac0: i32 = add 0x58109c0, 0x5813640 [ORD=113] [ID=8]<br>
0x58109c0: i32 = PTXISD::COPY_ADDRESS 0x5813540 [ID=7]<br>
0x5813540: i32 = TargetGlobalAddress<[16 x i8] addrspace(4)*<br>
@kernel_function.localArray> 0 [ID=4]<br>
0x5813640: i32,ch = load 0x57fa148, 0x5810dc0,<br>
0x58105c0<LD4[%retval.i]> [ORD=110] [ID=5]<br>
0x5810dc0: i32 = FrameIndex<0> [ORD=110] [ID=1]<br>
0x58105c0: i32 = undef [ORD=110] [ID=2]<br>
0x58105c0: i32 = undef [ORD=110] [ID=2]<br>
<br>
The command I am using is:<br>
<br>
clang kernels/<a href="http://fatal_error_test.cl" target="_blank">fatal_error_test.cl</a> -O0 -include ocldef.h -include<br>
<a href="http://builtin_functions_ptx.cl" target="_blank">builtin_functions_ptx.cl</a><br>
-D__x86_64__<br>
-ccc-host-triple ptx32 -Xclang<br>
-target-feature<br>
-Xclang +ptx23 -Xclang<br>
-target-feature<br>
-Xclang +compute20<br>
<br>
Any ideas ?<u></u><u></u></p>
<div>
<p class="MsoNormal"><u></u> <u></u></p>
</div>
<div>
<p class="MsoNormal">Unfortunately, this sample will not work at this time. First, the backend does not support i8 types yet. Second, at higher optimization levels, LLVM turns this loop into a memset intrinsic, which is also not yet implemented. :(<u></u><u></u></p>
</div>
<div>
<p class="MsoNormal"><u></u> <u></u></p>
</div>
<div>
<p class="MsoNormal">Hopefully I'll get some time soon to work on this, and other deficiencies. Patches are always welcome, too.<u></u><u></u></p>
</div>
<div>
<p class="MsoNormal"> <u></u><u></u></p>
</div>
<blockquote style="border:none;border-left:solid #CCCCCC 1.0pt;padding:0in 0in 0in 6.0pt;margin-left:4.8pt;margin-right:0in">
<p class="MsoNormal"><br>
Best regards<br>
<br>
Alberto<br>
_______________________________________________<br>
LLVM Developers mailing list<br>
<a href="mailto:LLVMdev@cs.uiuc.edu" target="_blank">LLVMdev@cs.uiuc.edu</a> <a href="http://llvm.cs.uiuc.edu" target="_blank">
http://llvm.cs.uiuc.edu</a><br>
<a href="http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev" target="_blank">http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev</a><u></u><u></u></p>
</blockquote>
</div>
<p class="MsoNormal"><br>
<br clear="all">
<u></u><u></u></p>
<div>
<p class="MsoNormal"><u></u> <u></u></p>
</div>
<p class="MsoNormal" style="margin-bottom:12.0pt">-- <u></u><u></u></p>
<div>
<p class="MsoNormal">Thanks,<u></u><u></u></p>
</div>
<div>
<p class="MsoNormal"><u></u> <u></u></p>
</div>
<div>
<p class="MsoNormal">Justin Holewinski<u></u><u></u></p>
</div>
<p class="MsoNormal"><u></u> <u></u></p>
</div></div></div>
</div>
</div>
</blockquote></div><br><br clear="all"><div><br></div>-- <br><br><div>Thanks,</div><div><br></div><div>Justin Holewinski</div><br>