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

broadcom/compiler: handle 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:54:45 +02:00
committed by Marge Bot
parent 9bf0b3a112
commit 742984a325
2 changed files with 15 additions and 3 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/broadcom/compiler/nir_to_vir.c
+8 -3
View File
@@ -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,
src/broadcom/compiler/v3d_compiler.h
+7
View File
@@ -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.
*