fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2026-04-19 16:40:36 +02:00
Code Issues Projects Releases Packages Wiki Activity Actions 1

v3d: support variable shared memory

Browse source 
Reviewed-by: Iago Toral Quiroga <itoral@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/25362>
This commit is contained in:
Karol Herbst 2023-09-23 11:58:54 +02:00 • committed by Marge Bot
parent 61b1a14e91
commit 6768c81974
3 changed files with 11 additions and 3 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

1
src/gallium/drivers/v3d/v3d_context.h
View file

@ -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;

5
src/gallium/drivers/v3d/v3d_uniforms.c
View file

@ -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;

8
src/gallium/drivers/v3d/v3dx_draw.c
View file

@ -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) {