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.ves...@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.ves...@rutgers.edu>
signature.asc
Description: This is a digitally signed message part
_______________________________________________ mesa-dev mailing list mesa-dev@lists.freedesktop.org https://lists.freedesktop.org/mailman/listinfo/mesa-dev
