AlexIndustrial/mesa
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

gallium: add PIPE_COMPUTE_CAP_MAX_SUBGROUPS

Browse Source 
This is required for enabling OpenCL subgroups.

Signed-off-by: Karol Herbst <git@karolherbst.de>
Reviewed-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/22893>
This commit is contained in:
Karol Herbst
2023-05-07 01:44:35 +02:00
committed by Marge Bot
parent 98d6edc6ad
commit f0b2944a0b
15 changed files with 49 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files
docs/gallium/screen.rst
+4
View File
@@ -798,6 +798,10 @@ pipe_screen::get_compute_param.
clock in MHz. Value type: ``uint32_t``
* ``PIPE_COMPUTE_CAP_MAX_COMPUTE_UNITS``: Maximum number of compute units
Value type: ``uint32_t``
* ``PIPE_COMPUTE_CAP_MAX_SUBGROUPS``: The max amount of subgroups there can be
inside a block. Non 0 indicates support for OpenCL subgroups including
implementing ``get_compute_state_subgroup_size`` if multiple subgroup sizes
are supported.
* ``PIPE_COMPUTE_CAP_IMAGES_SUPPORTED``: Whether images are supported
non-zero means yes, zero means no. Value type: ``uint32_t``
* ``PIPE_COMPUTE_CAP_SUBGROUP_SIZE``: The size of a basic execution unit in
src/gallium/drivers/asahi/agx_pipe.c
+3
View File
@@ -1858,6 +1858,9 @@ agx_get_compute_param(struct pipe_screen *pscreen, enum pipe_shader_ir ir_type,
case PIPE_COMPUTE_CAP_SUBGROUP_SIZE:
RET((uint32_t[]){32});
case PIPE_COMPUTE_CAP_MAX_SUBGROUPS:
RET((uint32_t[]){0 /* TODO */});
case PIPE_COMPUTE_CAP_MAX_VARIABLE_THREADS_PER_BLOCK:
RET((uint64_t[]){1024}); // TODO
}
src/gallium/drivers/crocus/crocus_screen.c
+1
View File
@@ -601,6 +601,7 @@ crocus_get_compute_param(struct pipe_screen *pscreen,
case PIPE_COMPUTE_CAP_MAX_GLOBAL_SIZE:
case PIPE_COMPUTE_CAP_MAX_PRIVATE_SIZE:
case PIPE_COMPUTE_CAP_MAX_INPUT_SIZE:
case PIPE_COMPUTE_CAP_MAX_SUBGROUPS:
// XXX: I think these are for Clover...
return 0;
src/gallium/drivers/freedreno/freedreno_screen.c
+3
View File
@@ -846,6 +846,9 @@ fd_get_compute_param(struct pipe_screen *pscreen, enum pipe_shader_ir ir_type,
case PIPE_COMPUTE_CAP_SUBGROUP_SIZE:
RET((uint32_t[]){32}); // TODO
case PIPE_COMPUTE_CAP_MAX_SUBGROUPS:
RET((uint32_t[]){0}); // TODO
case PIPE_COMPUTE_CAP_MAX_VARIABLE_THREADS_PER_BLOCK:
RET((uint64_t[]){ compiler->max_variable_workgroup_size });
}
src/gallium/drivers/iris/iris_screen.c
+3
View File
@@ -620,6 +620,9 @@ iris_get_compute_param(struct pipe_screen *pscreen,
case PIPE_COMPUTE_CAP_SUBGROUP_SIZE:
RET((uint32_t []) { BRW_SUBGROUP_SIZE });
case PIPE_COMPUTE_CAP_MAX_SUBGROUPS:
RET((uint32_t []) { devinfo->max_cs_workgroup_threads });
case PIPE_COMPUTE_CAP_MAX_MEM_ALLOC_SIZE:
case PIPE_COMPUTE_CAP_MAX_GLOBAL_SIZE:
RET((uint64_t []) { 1 << 30 }); /* TODO */
src/gallium/drivers/llvmpipe/lp_screen.c
+6
View File
@@ -537,6 +537,12 @@ llvmpipe_get_compute_param(struct pipe_screen *_screen,
*subgroup_size = lp_native_vector_width / 32;
}
return sizeof(uint32_t);
case PIPE_COMPUTE_CAP_MAX_SUBGROUPS:
if (ret) {
uint32_t *subgroup_size = ret;
*subgroup_size = 1024 / (lp_native_vector_width / 32);
}
return sizeof(uint32_t);
case PIPE_COMPUTE_CAP_MAX_COMPUTE_UNITS:
if (ret) {
uint32_t *max_compute_units = ret;
src/gallium/drivers/nouveau/nv50/nv50_screen.c
+2
View File
@@ -451,6 +451,8 @@ nv50_screen_get_compute_param(struct pipe_screen *pscreen,
RET((uint64_t []) { 4096 });
case PIPE_COMPUTE_CAP_SUBGROUP_SIZE:
RET((uint32_t []) { 32 });
case PIPE_COMPUTE_CAP_MAX_SUBGROUPS:
RET((uint32_t []) { 0 });
case PIPE_COMPUTE_CAP_MAX_MEM_ALLOC_SIZE:
RET((uint64_t []) { nouveau_device_get_global_mem_size(dev) });
case PIPE_COMPUTE_CAP_IMAGES_SUPPORTED:
src/gallium/drivers/nouveau/nvc0/nvc0_screen.c
+2
View File
@@ -558,6 +558,8 @@ nvc0_screen_get_compute_param(struct pipe_screen *pscreen,
RET((uint64_t []) { 4096 });
case PIPE_COMPUTE_CAP_SUBGROUP_SIZE:
RET((uint32_t []) { 32 });
case PIPE_COMPUTE_CAP_MAX_SUBGROUPS:
RET((uint32_t []) { 0 });
case PIPE_COMPUTE_CAP_MAX_MEM_ALLOC_SIZE:
RET((uint64_t []) { nouveau_device_get_global_mem_size(dev) });
case PIPE_COMPUTE_CAP_IMAGES_SUPPORTED:
src/gallium/drivers/panfrost/pan_screen.c
+3
View File
@@ -754,6 +754,9 @@ panfrost_get_compute_param(struct pipe_screen *pscreen,
case PIPE_COMPUTE_CAP_SUBGROUP_SIZE:
RET((uint32_t[]){pan_subgroup_size(dev->arch)});
case PIPE_COMPUTE_CAP_MAX_SUBGROUPS:
RET((uint32_t[]){0 /* TODO */});
case PIPE_COMPUTE_CAP_MAX_VARIABLE_THREADS_PER_BLOCK:
RET((uint64_t[]){1024}); // TODO
}
src/gallium/drivers/r600/r600_pipe_common.c
+1
View File
@@ -1033,6 +1033,7 @@ static int r600_get_compute_param(struct pipe_screen *screen,
}
return sizeof(uint32_t);
case PIPE_COMPUTE_CAP_MAX_PRIVATE_SIZE:
case PIPE_COMPUTE_CAP_MAX_SUBGROUPS:
break; /* unused */
case PIPE_COMPUTE_CAP_SUBGROUP_SIZE:
if (ret) {
src/gallium/drivers/radeonsi/si_get.c
+15
View File
@@ -1102,6 +1102,21 @@ static int si_get_compute_param(struct pipe_screen *screen, enum pipe_shader_ir
return sizeof(uint32_t);
case PIPE_COMPUTE_CAP_MAX_PRIVATE_SIZE:
break; /* unused */
case PIPE_COMPUTE_CAP_MAX_SUBGROUPS: {
if (ret) {
uint32_t *max_subgroups = ret;
unsigned threads = get_max_threads_per_block(sscreen, ir_type);
unsigned subgroup_size;
if (sscreen->debug_flags & DBG(W64_CS) || sscreen->info.gfx_level < GFX10)
subgroup_size = 64;
else
subgroup_size = 32;
*max_subgroups = threads / subgroup_size;
}
return sizeof(uint32_t);
}
case PIPE_COMPUTE_CAP_SUBGROUP_SIZE:
if (ret) {
uint32_t *subgroup_size = ret;
src/gallium/drivers/softpipe/sp_screen.c
+1
View File
@@ -557,6 +557,7 @@ softpipe_get_compute_param(struct pipe_screen *_screen,
case PIPE_COMPUTE_CAP_MAX_COMPUTE_UNITS:
case PIPE_COMPUTE_CAP_IMAGES_SUPPORTED:
case PIPE_COMPUTE_CAP_SUBGROUP_SIZE:
case PIPE_COMPUTE_CAP_MAX_SUBGROUPS:
case PIPE_COMPUTE_CAP_ADDRESS_BITS:
case PIPE_COMPUTE_CAP_MAX_VARIABLE_THREADS_PER_BLOCK:
break;
src/gallium/drivers/v3d/v3d_screen.c
+3
View File
@@ -548,6 +548,9 @@ v3d_get_compute_param(struct pipe_screen *pscreen, enum pipe_shader_ir ir_type,
case PIPE_COMPUTE_CAP_SUBGROUP_SIZE:
RET((uint32_t []) { 16 });
case PIPE_COMPUTE_CAP_MAX_SUBGROUPS:
RET((uint32_t []) { 0 });
}
return 0;
src/gallium/drivers/zink/zink_screen.c
+1
View File
@@ -440,6 +440,7 @@ zink_get_compute_param(struct pipe_screen *pscreen, enum pipe_shader_ir ir_type,
case PIPE_COMPUTE_CAP_SUBGROUP_SIZE:
RET((uint32_t []) { screen->info.props11.subgroupSize });
case PIPE_COMPUTE_CAP_MAX_SUBGROUPS:
case PIPE_COMPUTE_CAP_MAX_MEM_ALLOC_SIZE:
case PIPE_COMPUTE_CAP_MAX_CLOCK_FREQUENCY:
case PIPE_COMPUTE_CAP_MAX_COMPUTE_UNITS:
src/gallium/include/pipe/p_defines.h
+1
View File
@@ -1139,6 +1139,7 @@ enum pipe_compute_cap
PIPE_COMPUTE_CAP_MAX_MEM_ALLOC_SIZE,
PIPE_COMPUTE_CAP_MAX_CLOCK_FREQUENCY,
PIPE_COMPUTE_CAP_MAX_COMPUTE_UNITS,
PIPE_COMPUTE_CAP_MAX_SUBGROUPS,
PIPE_COMPUTE_CAP_IMAGES_SUPPORTED,
PIPE_COMPUTE_CAP_SUBGROUP_SIZE,
PIPE_COMPUTE_CAP_MAX_VARIABLE_THREADS_PER_BLOCK,