fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2025-12-25 19:30:11 +01:00
Code Issues Projects Releases Packages Wiki Activity Actions

panfrost: fix cl_local_size for precompiled shaders

Browse source 
nir_lower_compute_system_values will attempt to lower
load_workgroup_size unless workgroup_size_variable is set. For precomp
shaders, the workgroup size is set statically for each entrypoint by
nir_precompiled_build_variant. Because we call
lower_compute_system_values early, it sets the workgroup size to zero.
Temporarily setting workgroup_size_variable while we are still
processing all the entrypoints together inhibits this.

Signed-off-by: Olivia Lee <olivia.lee@collabora.com>
Fixes: 20970bcd96 ("panfrost: Add base of OpenCL C infrastructure")
Reviewed-by: Eric R. Smith <eric.smith@collabora.com>
Reviewed-by: Mary Guillemard <mary.guillemard@collabora.com>
Acked-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
(cherry picked from commit a410d90fd2)

Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38010>
This commit is contained in:
Olivia Lee 2025-10-08 21:57:00 -07:00 • committed by Dylan Baker
parent 10475e8ac1
commit 4202ea6c7f
3 changed files with 8 additions and 1 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

2
.pick_status.json
View file

@ -164,7 +164,7 @@
"description": "panfrost: fix cl_local_size for precompiled shaders",
"nominated": true,
"nomination_type": 2,
"resolution": 0,
"resolution": 1,
"main_sha": null,
"because_sha": "20970bcd9652ea8727a307ced9b372fad4b89991",
"notes": null

1
src/compiler/nir/nir_precompiled.h
View file

@ -652,6 +652,7 @@ nir_precompiled_build_variant(const nir_function *libfunc,
assert(libfunc->workgroup_size[0] != 0 && "must set workgroup size");
b.shader->info.workgroup_size_variable = false;
b.shader->info.workgroup_size[0] = libfunc->workgroup_size[0];
b.shader->info.workgroup_size[1] = libfunc->workgroup_size[1];
b.shader->info.workgroup_size[2] = libfunc->workgroup_size[2];

6
src/panfrost/clc/pan_compile.c
View file

@ -100,6 +100,12 @@ compile(void *memctx, const uint32_t *spirv, size_t spirv_size, unsigned arch)
nir_shader *nir =
spirv_to_nir(spirv, spirv_size / 4, NULL, 0, MESA_SHADER_KERNEL,
"library", &spirv_options, nir_options);
/* Workgroup size may be different between different entrypoints, so we
* mark it as variable to prevent it from being lowered to a constant while
* we are still processing all entrypoints together. This is tempoary,
* nir_precompiled_build_variant will set the fixed workgroup size for each
* entrypoint and set workgroup_size_variable back to false. */
nir->info.workgroup_size_variable = true;
nir_validate_shader(nir, "after spirv_to_nir");
nir_validate_ssa_dominance(nir, "after spirv_to_nir");
ralloc_steal(memctx, nir);