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

[Samuel Pitoiset]

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.

>
>   -ilia
>

-- 
-Samuel