From 4202ea6c7f8c437a48716464da23abf628d76d7e Mon Sep 17 00:00:00 2001 From: Olivia Lee Date: Wed, 8 Oct 2025 21:57:00 -0700 Subject: [PATCH] panfrost: fix cl_local_size for precompiled shaders 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 Fixes: 20970bcd965 ("panfrost: Add base of OpenCL C infrastructure") Reviewed-by: Eric R. Smith Reviewed-by: Mary Guillemard Acked-by: Alyssa Rosenzweig (cherry picked from commit a410d90fd254dd4930e72547de430a8cf5daf13b) Part-of: --- .pick_status.json | 2 +- src/compiler/nir/nir_precompiled.h | 1 + src/panfrost/clc/pan_compile.c | 6 ++++++ 3 files changed, 8 insertions(+), 1 deletion(-) diff --git a/.pick_status.json b/.pick_status.json index f695ec42d57..6cf4f41891e 100644 --- a/.pick_status.json +++ b/.pick_status.json @@ -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 diff --git a/src/compiler/nir/nir_precompiled.h b/src/compiler/nir/nir_precompiled.h index c94d70a1284..854c1f255e6 100644 --- a/src/compiler/nir/nir_precompiled.h +++ b/src/compiler/nir/nir_precompiled.h @@ -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]; diff --git a/src/panfrost/clc/pan_compile.c b/src/panfrost/clc/pan_compile.c index ef45a7a6314..63d78698c22 100644 --- a/src/panfrost/clc/pan_compile.c +++ b/src/panfrost/clc/pan_compile.c @@ -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);