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

[Jan Vesely]

On Wed, 2016-06-29 at 16:18 -0400, Jan Vesely wrote:
> 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)

atm, the only other info that is appended this way if global_offset
(for get_global_offset and get_global_id), the idea is to pass all
information this way and remove per driver upload of common info.

Jan

> 
> 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/7dc24333/attachment.sig>