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 <hdego...@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 <hdego...@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
_______________________________________________ mesa-dev mailing list mesa-dev@lists.freedesktop.org https://lists.freedesktop.org/mailman/listinfo/mesa-dev
