diff --git a/.pick_status.json b/.pick_status.json index 463b2b50cd3..34a26d9da04 100644 --- a/.pick_status.json +++ b/.pick_status.json @@ -44,7 +44,7 @@ "description": "rusticl/kernel: call nir_lower_variable_initializers earlier", "nominated": true, "nomination_type": 1, - "resolution": 0, + "resolution": 1, "main_sha": null, "because_sha": null, "notes": null diff --git a/src/gallium/frontends/rusticl/core/kernel.rs b/src/gallium/frontends/rusticl/core/kernel.rs index 6a3452400d5..9c9f8fb27dd 100644 --- a/src/gallium/frontends/rusticl/core/kernel.rs +++ b/src/gallium/frontends/rusticl/core/kernel.rs @@ -639,6 +639,12 @@ fn compile_nir_to_args( nir_pass!(nir, nir_scale_fdiv); nir.set_workgroup_size_variable_if_zero(); nir.structurize(); + nir_pass!( + nir, + nir_lower_variable_initializers, + nir_variable_mode::nir_var_function_temp + ); + while { let mut progress = false; nir_pass!(nir, nir_split_var_copies); diff --git a/src/gallium/frontends/rusticl/mesa/compiler/nir.rs b/src/gallium/frontends/rusticl/mesa/compiler/nir.rs index 207153025f4..15922386804 100644 --- a/src/gallium/frontends/rusticl/mesa/compiler/nir.rs +++ b/src/gallium/frontends/rusticl/mesa/compiler/nir.rs @@ -259,11 +259,6 @@ impl NirShader { } pub fn inline(&mut self, libclc: &NirShader) { - nir_pass!( - self, - nir_lower_variable_initializers, - nir_variable_mode::nir_var_function_temp, - ); nir_pass!(self, nir_lower_returns); nir_pass!(self, nir_link_shader_functions, libclc.nir.as_ptr()); nir_pass!(self, nir_inline_functions);