[Mesa-dev] [RFC 1/3] tgsi: Add WORK_DIM System Value

[Pierre Moreau]

On 05:43 PM - Apr 27 2016, Samuel Pitoiset wrote:
> 
> 
> On 04/27/2016 05:24 PM, Ilia Mirkin wrote:
> >On Wed, Apr 27, 2016 at 11:19 AM, Hans de Goede <hdegoede at 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.
> 
> I think your solution is fine too. Storing the number of dimensions passed
> to clEnqueueNDRangeKernel() as well as the local/global sizes makes sense.

Later versions of OpenCL are adding a few more of those (see [OpenCL 1.0][1]
compared to [OpenCL 2.0][2]). So putting it in a uniform would be better, but
we can just switch to a uniform the day we reach OpenCL 2.0.

Pierre


[1]: https://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/workItemFunctions.html
[2]: https://www.khronos.org/registry/cl/sdk/2.0/docs/man/xhtml/workItemFunctions.html

> 
> >
> >  -ilia
> >
> 
> -- 
> -Samuel
> _______________________________________________
> mesa-dev mailing list
> mesa-dev at lists.freedesktop.org
> https://lists.freedesktop.org/mailman/listinfo/mesa-dev
-------------- next part --------------
A non-text attachment was scrubbed...
Name: signature.asc
Type: application/pgp-signature
Size: 819 bytes
Desc: not available
URL: <https://lists.freedesktop.org/archives/mesa-dev/attachments/20160428/9ee2256e/attachment.sig>