microsoft/compiler: Move workgroup_size lowering from clc

It doesn't depend on the clc data being provided externally, so no
need to tie it there, we can re-use it for GL and Vulkan compute.

Reviewed-by: Sil Vilerino <sivileri@microsoft.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/14367>

src/microsoft/clc/clc_nir.c

@@ -67,22 +67,6 @@ lower_load_work_dim(nir_builder *b, nir_intrinsic_instr *intr,
    return true;
 }
 
-static bool
-lower_load_local_group_size(nir_builder *b, nir_intrinsic_instr *intr)
-{
-   b->cursor = nir_after_instr(&intr->instr);
-
-   nir_const_value v[3] = {
-      nir_const_value_for_int(b->shader->info.workgroup_size[0], 32),
-      nir_const_value_for_int(b->shader->info.workgroup_size[1], 32),
-      nir_const_value_for_int(b->shader->info.workgroup_size[2], 32)
-   };
-   nir_ssa_def *size = nir_build_imm(b, 3, 32, v);
-   nir_ssa_def_rewrite_uses(&intr->dest.ssa, size);
-   nir_instr_remove(&intr->instr);
-   return true;
-}
-
 static bool
 lower_load_num_workgroups(nir_builder *b, nir_intrinsic_instr *intr,
                           nir_variable *var)
@@ -146,9 +130,6 @@ clc_nir_lower_system_values(nir_shader *nir, nir_variable *var)
             case nir_intrinsic_load_work_dim:
                progress |= lower_load_work_dim(&b, intr, var);
                break;
-            case nir_intrinsic_load_workgroup_size:
-               lower_load_local_group_size(&b, intr);
-               break;
             case nir_intrinsic_load_num_workgroups:
                lower_load_num_workgroups(&b, intr, var);
                break;

src/microsoft/compiler/dxil_nir.c

@@ -1377,6 +1377,43 @@ dxil_nir_lower_system_values_to_zero(nir_shader* shader,
       &state);
 }
 
+static void
+lower_load_local_group_size(nir_builder *b, nir_intrinsic_instr *intr)
+{
+   b->cursor = nir_after_instr(&intr->instr);
+
+   nir_const_value v[3] = {
+      nir_const_value_for_int(b->shader->info.workgroup_size[0], 32),
+      nir_const_value_for_int(b->shader->info.workgroup_size[1], 32),
+      nir_const_value_for_int(b->shader->info.workgroup_size[2], 32)
+   };
+   nir_ssa_def *size = nir_build_imm(b, 3, 32, v);
+   nir_ssa_def_rewrite_uses(&intr->dest.ssa, size);
+   nir_instr_remove(&intr->instr);
+}
+
+static bool
+lower_system_values_impl(nir_builder *b, nir_instr *instr, void *_state)
+{
+   if (instr->type != nir_instr_type_intrinsic)
+      return false;
+   nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
+   switch (intr->intrinsic) {
+   case nir_intrinsic_load_workgroup_size:
+      lower_load_local_group_size(b, intr);
+      return true;
+   default:
+      return false;
+   }
+}
+
+bool
+dxil_nir_lower_system_values(nir_shader *shader)
+{
+   return nir_shader_instructions_pass(shader, lower_system_values_impl,
+      nir_metadata_block_index | nir_metadata_dominance | nir_metadata_loop_analysis, NULL);
+}
+
 static const struct glsl_type *
 get_bare_samplers_for_type(const struct glsl_type *type, bool is_shadow)
 {

src/microsoft/compiler/dxil_nir.h

@@ -48,6 +48,7 @@ bool dxil_nir_lower_double_math(nir_shader *shader);
 bool dxil_nir_lower_system_values_to_zero(nir_shader *shader,
                                           gl_system_value* system_value,
                                           uint32_t count);
+bool dxil_nir_lower_system_values(nir_shader *shader);
 bool dxil_nir_split_typed_samplers(nir_shader *shader);
 bool dxil_nir_lower_bool_input(struct nir_shader *s);
 bool dxil_nir_lower_sysval_to_load_input(nir_shader *s, nir_variable **sysval_vars);

src/microsoft/compiler/nir_to_dxil.c

@@ -4984,6 +4984,7 @@ nir_to_dxil(struct nir_shader *s, const struct nir_to_dxil_options *opts,
    NIR_PASS_V(s, nir_lower_frexp);
    NIR_PASS_V(s, nir_lower_flrp, 16 | 32 | 64, true);
    NIR_PASS_V(s, nir_lower_io, nir_var_shader_in | nir_var_shader_out, type_size_vec4, (nir_lower_io_options)0);
+   NIR_PASS_V(s, dxil_nir_lower_system_values);
 
    optimize_nir(s, opts);