From 2038134efc279f7ded143cf9c48b7d7a496eb2a7 Mon Sep 17 00:00:00 2001 From: Qiang Yu Date: Sun, 27 Apr 2025 10:47:11 +0800 Subject: [PATCH] radeonsi: calc workgroup size for mesh shader MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Reviewed-by: Marek Olšák Part-of: --- src/gallium/drivers/radeonsi/si_shader.c | 19 ++++++++++++++++--- 1 file changed, 16 insertions(+), 3 deletions(-) diff --git a/src/gallium/drivers/radeonsi/si_shader.c b/src/gallium/drivers/radeonsi/si_shader.c index 709fa42122e..fb7537c6ba4 100644 --- a/src/gallium/drivers/radeonsi/si_shader.c +++ b/src/gallium/drivers/radeonsi/si_shader.c @@ -117,6 +117,7 @@ unsigned si_shader_io_get_unique_index(unsigned semantic) unsigned si_get_max_workgroup_size(const struct si_shader *shader) { + struct si_screen *sscreen = shader->selector->screen; mesa_shader_stage stage = shader->is_gs_copy_shader ? MESA_SHADER_VERTEX : shader->selector->stage; @@ -130,18 +131,20 @@ unsigned si_get_max_workgroup_size(const struct si_shader *shader) return shader->info.num_streamout_vec4s ? 256 : 128; /* As part of merged shader. */ - return shader->selector->screen->info.gfx_level >= GFX9 && + return sscreen->info.gfx_level >= GFX9 && (shader->key.ge.as_ls || shader->key.ge.as_es) ? 128 : shader->wave_size; case MESA_SHADER_TESS_CTRL: /* Return this so that LLVM doesn't remove s_barrier * instructions on chips where we use s_barrier. */ - return shader->selector->screen->info.gfx_level >= GFX7 ? 128 : shader->wave_size; + return sscreen->info.gfx_level >= GFX7 ? 128 : shader->wave_size; case MESA_SHADER_GEOMETRY: /* GS can always generate up to 256 vertices. */ - return shader->selector->screen->info.gfx_level >= GFX9 ? 256 : shader->wave_size; + return sscreen->info.gfx_level >= GFX9 ? 256 : shader->wave_size; + case MESA_SHADER_TASK: + case MESA_SHADER_MESH: case MESA_SHADER_COMPUTE: break; /* see below */ @@ -157,6 +160,16 @@ unsigned si_get_max_workgroup_size(const struct si_shader *shader) unsigned max_work_group_size = (uint32_t)local_size[0] * (uint32_t)local_size[1] * (uint32_t)local_size[2]; + + /* Without multi-row export, we need at least number of output vertex/primitive + * threads in workgroup for export (one vertex/primitive per thread). + */ + if (stage == MESA_SHADER_MESH && !sscreen->info.mesh_fast_launch_2) { + max_work_group_size = MAX3(max_work_group_size, + shader->selector->info.base.mesh.max_vertices_out, + shader->selector->info.base.mesh.max_primitives_out); + } + assert(max_work_group_size); return max_work_group_size; }