From 742984a3256e11ff30eb84d51286f8743615b1ac Mon Sep 17 00:00:00 2001 From: Karol Herbst Date: Sat, 23 Sep 2023 11:54:45 +0200 Subject: [PATCH] broadcom/compiler: handle variable shared memory Reviewed-by: Iago Toral Quiroga Part-of: --- src/broadcom/compiler/nir_to_vir.c | 11 ++++++++--- src/broadcom/compiler/v3d_compiler.h | 7 +++++++ 2 files changed, 15 insertions(+), 3 deletions(-) diff --git a/src/broadcom/compiler/nir_to_vir.c b/src/broadcom/compiler/nir_to_vir.c index 634c8a17021..1c665cd18a6 100644 --- a/src/broadcom/compiler/nir_to_vir.c +++ b/src/broadcom/compiler/nir_to_vir.c @@ -4674,7 +4674,7 @@ nir_to_vir(struct v3d_compile *c) ffs(util_next_power_of_two(MAX2(wg_size, 64))) - 1; assert(c->local_invocation_index_bits <= 8); - if (c->s->info.shared_size) { + if (c->s->info.shared_size || c->s->info.cs.has_variable_shared_mem) { struct qreg wg_in_mem = vir_SHR(c, c->cs_payload[1], vir_uniform_ui(c, 16)); if (c->s->info.workgroup_size[0] != 1 || @@ -4686,8 +4686,13 @@ nir_to_vir(struct v3d_compile *c) wg_in_mem = vir_AND(c, wg_in_mem, vir_uniform_ui(c, wg_mask)); } - struct qreg shared_per_wg = - vir_uniform_ui(c, c->s->info.shared_size); + + struct qreg shared_per_wg; + if (c->s->info.cs.has_variable_shared_mem) { + shared_per_wg = vir_uniform(c, QUNIFORM_SHARED_SIZE, 0); + } else { + shared_per_wg = vir_uniform_ui(c, c->s->info.shared_size); + } c->cs_shared_offset = vir_ADD(c, diff --git a/src/broadcom/compiler/v3d_compiler.h b/src/broadcom/compiler/v3d_compiler.h index c40562bcc0e..0adbab9a45d 100644 --- a/src/broadcom/compiler/v3d_compiler.h +++ b/src/broadcom/compiler/v3d_compiler.h @@ -328,6 +328,13 @@ enum quniform_contents { */ QUNIFORM_SHARED_OFFSET, + /** + * OpenCL variable shared memory + * + * This will only be used when the shader declares variable_shared_memory. + */ + QUNIFORM_SHARED_SIZE, + /** * Returns the number of layers in the framebuffer. *