<div dir="ltr">I think the "<span style="font-size:12.8px">define linkonce_odr i32 @get_global_id(i32 %dim) #5" you dumped is</span><span style="font-size:12.8px"> llvm IR after inlining and opt.</span><div>The commit you mentioned(ba9858) doesn't change get_local_size() at all. <br></div><div><br></div><div>I never worked on OpenCL+HSA. I just wonder if libclc supports HSA.</div><div><br></div><div>HSA RT uses 'hsa_kernel_dispatch_packet_t' to get know workgroup size and grid size. so far, I didn't see hsa-specific implementation appears in libclc. </div><div><br></div><div>thanks,</div><div>--lx</div><div><br></div><div><br></div><div><br></div></div><div class="gmail_extra"><br><div class="gmail_quote">On Sun, Mar 6, 2016 at 1:28 AM, 李弘宇 <span dir="ltr"><<a href="mailto:zhenlinospirit@gmail.com" target="_blank">zhenlinospirit@gmail.com</a>></span> wrote:<br><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex"><div dir="ltr">Hi Mr. Liu, <div><br></div><div>Thanks for your quick reply.</div><div><br></div><div>I compiled the code with the libclc_trunk and linked the bitcode file under $LIBCLC_DIR/built_libs/tahiti-amdgcn--.bc. After looking into the libclc, it is currently using the new workitem intrinsics (commit ba9858caa1e927a6fcc601e3466faa693835db5e). In the linked bitcode ($LIBCLC_DIR/built_libs/tahiti-amdgcn--.bc), it has the following code segment,</div><div><br></div><div><div>define linkonce_odr i32 @get_global_id(i32 %dim) #5 {</div><div>entry:</div><div> switch i32 %dim, label %get_local_id.exit [</div><div> i32 0, label %get_group_id.exit.thread</div><div> i32 1, label %get_group_id.exit.thread22</div><div> i32 2, label %get_group_id.exit.thread24</div><div> ]</div><div><br></div><div>get_group_id.exit.thread: ; preds = %entry</div><div> %x.i = tail call i32 @llvm.amdgcn.workgroup.id.x() #13</div><div> %x.i12 = tail call i32 @llvm.r600.read.local.size.x() #3</div><div> %mul26 = mul i32 %x.i12, %x.i</div><div> %x.i4 = tail call i32 @llvm.amdgcn.workitem.id.x() #13, !range !1</div><div> br label %get_local_id.exit</div><div>...</div><div>}</div></div><div><br></div><div>So it shows that some intrinstics are still using llvm.r600.xxx. I have no idea if I ever missed something so that it doesn't work.</div><div><br></div><div>Thanks.</div><div><br></div><div>Best regards,</div><span class=""><div><br></div><div><div><div style="color:rgb(0,0,0);font-size:12.8px">李弘宇 (Li, Hong-Yu)</div><div style="color:rgb(0,0,0);font-size:12.8px">Department of Computer Science & Information Engineering</div><div style="color:rgb(0,0,0);font-size:12.8px">National Taiwan University</div></div></div></span></div><div class="HOEnZb"><div class="h5"><div class="gmail_extra"><br><div class="gmail_quote">On Sun, Mar 6, 2016 at 12:59 AM, Liu Xin <span dir="ltr"><<a href="mailto:navy.xliu@gmail.com" target="_blank">navy.xliu@gmail.com</a>></span> wrote:<br><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex"><div dir="ltr">Li, Hong-Yu, <div><br></div><div>it's because get_group_id() uses get_local_size</div><div><div>_CLC_DEF size_t get_global_id(uint dim) {</div><div> return get_group_id(dim)*get_local_size(dim) + get_local_id(dim);</div><div>}</div></div><div><br></div><div>in libclc/amdgcn, 'get_local_size' invokes r600-xxx intrinsics. I doubt that libclc ever supports hsa-runtime before.</div><div><br></div><div><br></div><div>thanks,</div><div>--lx</div><div><br></div></div><div class="gmail_extra"><br><div class="gmail_quote"><div><div>On Sun, Mar 6, 2016 at 12:11 AM, 李弘宇 via llvm-dev <span dir="ltr"><<a href="mailto:llvm-dev@lists.llvm.org" target="_blank">llvm-dev@lists.llvm.org</a>></span> wrote:<br></div></div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex"><div><div><div dir="ltr">Dear Developers,<div><br></div><div><span style="color:rgb(0,0,0);font-size:12.8px">I compiled a OpenCL kernel before (on Nov. last year) like</span><br></div><div><span style="color:rgb(0,0,0);font-size:12.8px"><br></span></div><div><span style="color:rgb(0,0,0);font-size:12.8px">__kernel void g(__global float* array)</span></div><div><span style="color:rgb(0,0,0);font-size:12.8px">{</span></div><div><span style="color:rgb(0,0,0);font-size:12.8px"> array[get_global_id(0)] = 1;</span></div><div><span style="color:rgb(0,0,0);font-size:12.8px">}</span></div><div><span style="color:rgb(0,0,0);font-size:12.8px"><br></span></div><div><span style="color:rgb(0,0,0);font-size:12.8px">with libclc, </span><span style="color:rgb(0,0,0);font-size:12.8px">which would originally use the instrinsics like </span><font color="#000000"><span style="font-size:12.8px">llvm.r600.read.local.size.x().</span></font></div><div><font color="#000000"><span style="font-size:12.8px"><br></span></font></div><div><font color="#000000"><span style="font-size:12.8px">I executed the generated object file with one version of the hsa-runtime [1] provided by Mr. Stellard, when there was more than one workgroup, the output of the program wasn't correct at that time. I guessed this might be because get_group_id() always returned 1 (not quite sure what was going on at that time).</span></font></div><div><font color="#000000"><span style="font-size:12.8px"><br></span></font></div><div><font color="#000000"><span style="font-size:12.8px">When I compile such cases using current llvm trunk, it uses a set of instrinsics starting with llvm.amdgcn, while it still uses llvm.r600.read.local.size.x(). The output LLVM IR code is like:</span></font></div><div><font color="#000000"><span style="font-size:12.8px"><br></span></font></div><div><font color="#000000"><span style="font-size:12.8px">define void @g(float addrspace(1)* nocapture %array) #0 {</span><br></font></div><div><font color="#000000"><span style="font-size:12.8px"> %x.i.i = tail call i32 @llvm.amdgcn.workgroup.id.x() #2</span></font></div><div><font color="#000000"><span style="font-size:12.8px"> %x.i12.i = tail call i32 @llvm.r600.read.local.size.x() #1</span></font></div><div><font color="#000000"><span style="font-size:12.8px"> %mul26.i = mul i32 %x.i12.i, %x.i.i </span></font></div><div><font color="#000000"><span style="font-size:12.8px"> %x.i4.i = tail call i32 @llvm.amdgcn.workitem.id.x() #2, !range !7</span></font><br></div><div><font color="#000000"><span style="font-size:12.8px"> %add.i = add i32 %x.i4.i, %mul26.i</span></font></div><div><font color="#000000"><span style="font-size:12.8px"> %0 = sext i32 %add.i to i64</span></font></div><div><font color="#000000"><span style="font-size:12.8px"> %arrayidx = getelementptr inbounds float, float addrspace(1)* %array, i64 %0</span></font></div><div><font color="#000000"><span style="font-size:12.8px"> store float 1.000000e+00, float addrspace(1)* %arrayidx, align 4, !tbaa !8</span></font></div><div><font color="#000000"><span style="font-size:12.8px"> ret void</span></font></div><div><font color="#000000"><span style="font-size:12.8px">}</span></font></div><div><font color="#000000"><span style="font-size:12.8px"><br></span></font></div><div><font color="#000000"><span style="font-size:12.8px">which cannot be handled by llc with the message "the non-hsa instrinsic with hsa target shown". </span></font></div><div><font color="#000000"><span style="font-size:12.8px"><br></span></font></div><div><font color="#000000"><span style="font-size:12.8px">After looking into the log (r259297)</span></font><span style="font-size:12.8px;color:rgb(0,0,0)">, m</span><font color="#000000"><span style="font-size:12.8px">y question is that is there other intrinsic that support this case when the target is amdgcn--amdhsa? In the log of </span></font><span style="color:rgb(0,0,0);font-size:12.8px">r259297, it states that </span><font color="#000000"><span style="font-size:12.8px">AMDGPUPromoteAlloca pass (a backend pass) will generate this intrinsic, but even when I just emit-llvm without going through llc, this intrinsic is still emitted.</span></font></div><div><font color="#000000"><span style="font-size:12.8px"><br></span></font></div><div><span style="font-size:12.8px;color:rgb(0,0,0)">[1] <a href="https://github.com/tstellarAMD/hsa-runtime" target="_blank">https://github.com/tstellarAMD/hsa-runtime</a></span><br></div><div><font color="#000000"><span style="font-size:12.8px"><br></span></font></div><div><div style="color:rgb(0,0,0);font-size:12.8px"><br></div><div style="color:rgb(0,0,0);font-size:12.8px">Regards,</div><div style="color:rgb(0,0,0);font-size:12.8px"><br></div><div style="color:rgb(0,0,0);font-size:12.8px"><div style="font-size:12.8px">李弘宇 (Li, Hong-Yu)</div><div style="font-size:12.8px">Department of Computer Science & Information Engineering</div><div style="font-size:12.8px">National Taiwan University</div></div></div></div>
<br></div></div>_______________________________________________<br>
LLVM Developers mailing list<br>
<a href="mailto:llvm-dev@lists.llvm.org" target="_blank">llvm-dev@lists.llvm.org</a><br>
<a href="http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev" rel="noreferrer" target="_blank">http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev</a><br>
<br></blockquote></div><br></div>
</blockquote></div><br></div>
</div></div></blockquote></div><br></div>