[Mesa-dev] [Bug 86326] clEnqueueNDRangeKernel global_work_offset ignored

[bugzilla-daemon at freedesktop.org]

https://bugs.freedesktop.org/show_bug.cgi?id=86326

--- Comment #5 from Tom Stellard <tstellar at gmail.com> ---
Comment on attachment 114337
  --> https://bugs.freedesktop.org/attachment.cgi?id=114337
libclc get_global_offset() and get_global_id() patch

Review of attachment 114337:
-----------------------------------------------------------------

Hi,

Thanks for the patch.  Would you be able to send your updated version to:
libclc-dev at pcc.me.uk

::: ptx-nvidiacl/lib/workitem/get_global_offset.cl
@@ +6,5 @@
> +  case 1:  return __builtin_ptx_read_global_offset_y();
> +  case 2:  return __builtin_ptx_read_global_offset_z();*/
> +  default: return 0;
> +  }
> +}

Why is this commented out?  Also the function name is wrong:
get_local_offset().

::: r600/lib/workitem/get_global_offset.ll
@@ +14,5 @@
> +  %z = call i32 @llvm.AMDGPU.read.global.offset.z() nounwind readnone
> +  ret i32 %z
> +default:
> +  ret i32 0
> +}

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 ...

-- 
You are receiving this mail because:
You are the assignee for the bug.
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.freedesktop.org/archives/mesa-dev/attachments/20150316/f5402bc5/attachment.html>