<html>
<head>
<base href="https://bugs.freedesktop.org/" />
</head>
<body>
<p>
<div>
<b><a class="bz_bug_link
bz_status_NEW "
title="NEW - clEnqueueNDRangeKernel global_work_offset ignored"
href="https://bugs.freedesktop.org/show_bug.cgi?id=86326#c5">Comment # 5</a>
on <a class="bz_bug_link
bz_status_NEW "
title="NEW - clEnqueueNDRangeKernel global_work_offset ignored"
href="https://bugs.freedesktop.org/show_bug.cgi?id=86326">bug 86326</a>
from <span class="vcard"><a class="email" href="mailto:tstellar@gmail.com" title="Tom Stellard <tstellar@gmail.com>"> <span class="fn">Tom Stellard</span></a>
</span></b>
<pre>Comment on <span class=""><a href="attachment.cgi?id=114337" name="attach_114337" title="libclc get_global_offset() and get_global_id() patch">attachment 114337</a> <a href="attachment.cgi?id=114337&action=edit" title="libclc get_global_offset() and get_global_id() patch">[details]</a></span> <a href='page.cgi?id=splinter.html&bug=86326&attachment=114337'>[review]</a>
libclc get_global_offset() and get_global_id() patch
Review of <span class=""><a href="attachment.cgi?id=114337" name="attach_114337" title="libclc get_global_offset() and get_global_id() patch">attachment 114337</a> <a href="attachment.cgi?id=114337&action=edit" title="libclc get_global_offset() and get_global_id() patch">[details]</a></span> <a href='page.cgi?id=splinter.html&bug=86326&attachment=114337'>[review]</a>:
-----------------------------------------------------------------
Hi,
Thanks for the patch. Would you be able to send your updated version to:
<a href="mailto:libclc-dev@pcc.me.uk">libclc-dev@pcc.me.uk</a>
::: ptx-nvidiacl/lib/workitem/get_global_offset.cl
@@ +6,5 @@
<span class="quote">> + case 1: return __builtin_ptx_read_global_offset_y();
> + case 2: return __builtin_ptx_read_global_offset_z();*/
> + default: return 0;
> + }
> +}</span >
Why is this commented out? Also the function name is wrong:
get_local_offset().
::: r600/lib/workitem/get_global_offset.ll
@@ +14,5 @@
<span class="quote">> + %z = call i32 @llvm.AMDGPU.read.global.offset.z() nounwind readnone
> + ret i32 %z
> +default:
> + ret i32 0
> +}</span >
This should be implemented in OpenCL C rather than LLVM IR. Now that
DataLayouts are mandatory, it makes implementing common code in LLVM IR much
more difficult.
You can use the builtins defined in your previous patch for this:
__builtin_amdgpu_read_global_offset_x ...</pre>
</div>
</p>
<hr>
<span>You are receiving this mail because:</span>
<ul>
<li>You are the assignee for the bug.</li>
</ul>
</body>
</html>