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