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

[Hans de Goede]

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.

Regards,

Hans



>
> Is this the equivalent of gl_LocalGroupSizeARB from
> GL_ARB_compute_variable_group_size ? [which is basically the same as
> gl_WorkGroupSize, but allowed to be specified at dispatch time]
>
>
> On Wed, Apr 27, 2016 at 10:43 AM, Hans de Goede <hdegoede at redhat.com> wrote:
>> Add a new WORK_DIM SV type, this is will return the grid dimensions
>> (1-4) for compute (opencl) kernels.
>>
>> This is necessary to implement the opencl get_work_dim() function.
>>
>> Signed-off-by: Hans de Goede <hdegoede at redhat.com>
>> ---
>>  src/gallium/auxiliary/tgsi/tgsi_strings.c  | 1 +
>>  src/gallium/include/pipe/p_shader_tokens.h | 1 +
>>  2 files changed, 2 insertions(+)
>>
>> diff --git a/src/gallium/auxiliary/tgsi/tgsi_strings.c b/src/gallium/auxiliary/tgsi/tgsi_strings.c
>> index 894d475..f65d7b4 100644
>> --- a/src/gallium/auxiliary/tgsi/tgsi_strings.c
>> +++ b/src/gallium/auxiliary/tgsi/tgsi_strings.c
>> @@ -100,6 +100,7 @@ const char *tgsi_semantic_names[TGSI_SEMANTIC_COUNT] =
>>     "HELPER_INVOCATION",
>>     "BASEINSTANCE",
>>     "DRAWID",
>> +   "WORK_DIM",
>>  };
>>
>>  const char *tgsi_texture_names[TGSI_TEXTURE_COUNT] =
>> diff --git a/src/gallium/include/pipe/p_shader_tokens.h b/src/gallium/include/pipe/p_shader_tokens.h
>> index 514b339..d8ded50 100644
>> --- a/src/gallium/include/pipe/p_shader_tokens.h
>> +++ b/src/gallium/include/pipe/p_shader_tokens.h
>> @@ -200,6 +200,7 @@ enum tgsi_semantic {
>>     TGSI_SEMANTIC_HELPER_INVOCATION,  /**< current invocation is helper */
>>     TGSI_SEMANTIC_BASEINSTANCE,
>>     TGSI_SEMANTIC_DRAWID,
>> +   TGSI_SEMANTIC_WORK_DIM,    /**< opencl get_work_dim value */
>>     TGSI_SEMANTIC_COUNT,       /**< number of semantic values */
>>  };
>>
>> --
>> 2.7.4
>>