[Mesa-dev] [PATCH v2 2/4] clover: Pass work_dim parameter of clEnqueueNDRangeKernel() to driver

[Francisco Jerez]

Ilia Mirkin <imirkin at alum.mit.edu> writes:

> On Wed, Jun 29, 2016 at 2:42 PM, Jan Vesely <jan.vesely at rutgers.edu> wrote:
>> On Wed, 2016-06-29 at 14:37 +0200, Hans de Goede wrote:
>>> In order to implement get_work_dim() the driver may need to know the
>>> clEnqueueNDRangeKernel() work_dim parameter, so pass it to the
>>> driver.
>>
>> The workdim info is passed as the first implicit argument (after
>> explicit kernel arguments, see invocation.cpp:600 ). Can you use that?
>> Both radeonsi and r600 already do.
>
> How does the driver know where the implicit arguments start?
>

The plan was to let the LLVM back-end specify where it wants them (see
XXX comment in clover/llvm/invocation.cpp:490) e.g. via some
clover-specific kernel argument metadata annotations telling the
front-end to initialize a given kernel argument implicitly, but nobody
implemented the changes required in the AMDGPU LLVM back-end so right
now the grid dimension and offset components are hard-coded to be passed
at the end of the input buffer in that order.

>   -ilia
> _______________________________________________
> mesa-dev mailing list
> mesa-dev at lists.freedesktop.org
> https://lists.freedesktop.org/mailman/listinfo/mesa-dev
-------------- next part --------------
A non-text attachment was scrubbed...
Name: signature.asc
Type: application/pgp-signature
Size: 212 bytes
Desc: not available
URL: <https://lists.freedesktop.org/archives/mesa-dev/attachments/20160629/efed07e5/attachment.sig>