On Wed, Apr 27, 2016 at 11:19 AM, Hans de Goede <hdego...@redhat.com> wrote: > Hi, > > On 27-04-16 16:49, Ilia Mirkin wrote: >> >> Please add this semantic to src/gallium/docs and explain what it >> means. > > > Ah, I was under the impression these were not documented, but I > now see they are, will fix. > >> (I'm not even sure what this is, and the easily-found opencl >> >> docs helpfully indicate that get_work_dim returns the work_dim...) >> This wasn't done for some of the other compute-specific semantics. > > > When launching a grid through clEnqueueNDRangeKernel the user > can tell how many dimensions of the pipe_grid_info->grid / > pipe_grid_info->block equivalent the user is actually passing in > this is done through a parameter to that function called work_dim: > > https://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/clEnqueueNDRangeKernel.html > > Clover pads the y and z values (if unused) to 1, so it always > ends up giving a 3 dimensional grid and block sizes to the driver. > > But an opencl kernel can use the opencl builtin get_work_dim > function, which returns how much dimensions where actually > passed into clEnqueueNDRangeKernel. > > Note that passing in e.g. work_dim = 3 with a grid size of { 4, 1, 1 } > is perfectly legal and then get_work_dim should actually return 3. > > So we need to store the clEnqueueNDRangeKernel workdim parameter > somewhere, and allow getting it from a kernel. Using a system-value > seems the best solution for this to me. > > I hope this helps to explain things, if not keep asking :) > > Also if you've a better solution, I'm all ears.
Sounds like there's no real way to distinguish it from the kernel, the parameters given to the hw are all the same, so it's just a thing that needs to be passed along. I think that what you have is a fine solution, although I question the usefulness of this being an OpenCL builtin. But we're not about to change that now. An alternative would be for clover to stick it into a uniform and have the kernel read it out from there. If there are going to be a ton of these, that might be preferable. If it's just the one (or maybe two), not worth it. -ilia _______________________________________________ mesa-dev mailing list mesa-dev@lists.freedesktop.org https://lists.freedesktop.org/mailman/listinfo/mesa-dev