Ilia Mirkin <imir...@alum.mit.edu> writes: > 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 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@lists.freedesktop.org > https://lists.freedesktop.org/mailman/listinfo/mesa-dev
signature.asc
Description: PGP signature
_______________________________________________ mesa-dev mailing list mesa-dev@lists.freedesktop.org https://lists.freedesktop.org/mailman/listinfo/mesa-dev
