<div class="gmail_quote">On Mon, Nov 14, 2011 at 8:57 AM, Alberto Magni <span dir="ltr"><<a href="mailto:alberto.magni86@gmail.com">alberto.magni86@gmail.com</a>></span> wrote:<br><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex;">
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 ?<br></blockquote><div><br></div><div>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. :(</div>
<div><br></div><div>Hopefully I'll get some time soon to work on this, and other deficiencies. Patches are always welcome, too.</div><div> </div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex;">
<br>
Best regards<br>
<br>
Alberto<br>
_______________________________________________<br>
LLVM Developers mailing list<br>
<a href="mailto:LLVMdev@cs.uiuc.edu">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><br>
</blockquote></div><br><br clear="all"><div><br></div>-- <br><br><div>Thanks,</div><div><br></div><div>Justin Holewinski</div><br>