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

[Jan Vesely]

On Wed, 2016-06-29 at 15:45 -0400, Ilia Mirkin wrote:
> 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 compiler knows the number and types of arguments, it can compute
the offset in the argument vector where implicit arguments start.
you can either have compiler intrinsic per value, or one intrinsic to
get implicitarg.pointer and load from there. (r600 and radeonsi are in
the process of switching to the latter)

Jan

> 
>   -ilia
-- 

Jan Vesely <jan.vesely at rutgers.edu>
-------------- next part --------------
A non-text attachment was scrubbed...
Name: signature.asc
Type: application/pgp-signature
Size: 819 bytes
Desc: This is a digitally signed message part
URL: <https://lists.freedesktop.org/archives/mesa-dev/attachments/20160629/8b283694/attachment.sig>