From 6768c8197486f49db07a8ce0a5d4b5bd365a8cfa Mon Sep 17 00:00:00 2001 From: Karol Herbst Date: Sat, 23 Sep 2023 11:58:54 +0200 Subject: [PATCH] v3d: support variable shared memory Reviewed-by: Iago Toral Quiroga Part-of: --- src/gallium/drivers/v3d/v3d_context.h | 1 + src/gallium/drivers/v3d/v3d_uniforms.c | 5 +++++ src/gallium/drivers/v3d/v3dx_draw.c | 8 +++++--- 3 files changed, 11 insertions(+), 3 deletions(-) diff --git a/src/gallium/drivers/v3d/v3d_context.h b/src/gallium/drivers/v3d/v3d_context.h index 966a809eb11..dc855408f62 100644 --- a/src/gallium/drivers/v3d/v3d_context.h +++ b/src/gallium/drivers/v3d/v3d_context.h @@ -577,6 +577,7 @@ struct v3d_context { uint32_t compute_num_workgroups[3]; uint32_t compute_workgroup_size[3]; struct v3d_bo *compute_shared_memory; + uint32_t shared_memory; struct v3d_vertex_stateobj *vtx; diff --git a/src/gallium/drivers/v3d/v3d_uniforms.c b/src/gallium/drivers/v3d/v3d_uniforms.c index 60beab84b3c..f9da74454ec 100644 --- a/src/gallium/drivers/v3d/v3d_uniforms.c +++ b/src/gallium/drivers/v3d/v3d_uniforms.c @@ -382,6 +382,10 @@ v3d_write_uniforms(struct v3d_context *v3d, struct v3d_job *job, v3d->compute_shared_memory, 0); break; + case QUNIFORM_SHARED_SIZE: + cl_aligned_u32(&uniforms, v3d->shared_memory); + break; + case QUNIFORM_FB_LAYERS: cl_aligned_u32(&uniforms, job->num_layers); break; @@ -471,6 +475,7 @@ v3d_set_shader_uniform_dirty_flags(struct v3d_compiled_shader *shader) case QUNIFORM_NUM_WORK_GROUPS: case QUNIFORM_WORK_GROUP_SIZE: case QUNIFORM_SHARED_OFFSET: + case QUNIFORM_SHARED_SIZE: /* Compute always recalculates uniforms. */ break; diff --git a/src/gallium/drivers/v3d/v3dx_draw.c b/src/gallium/drivers/v3d/v3dx_draw.c index c54ad4cc02f..ca378096e46 100644 --- a/src/gallium/drivers/v3d/v3dx_draw.c +++ b/src/gallium/drivers/v3d/v3dx_draw.c @@ -1526,12 +1526,14 @@ v3d_launch_grid(struct pipe_context *pctx, const struct pipe_grid_info *info) if (v3d->prog.compute->prog_data.base->threads == 4) submit.cfg[5] |= V3D_CSD_CFG5_THREADING; - if (v3d->prog.compute->prog_data.compute->shared_size) { + uint32_t shared_size = v3d->prog.compute->prog_data.compute->shared_size + + info->variable_shared_mem; + if (shared_size) { v3d->compute_shared_memory = v3d_bo_alloc(v3d->screen, - v3d->prog.compute->prog_data.compute->shared_size * - num_wgs, + shared_size * num_wgs, "shared_vars"); + v3d->shared_memory = shared_size; } util_dynarray_foreach(&v3d->global_buffers, struct pipe_resource *, res) {