nir: Fix lowering terminology of compute system values: "from"->"to".

This is to match other NIR terminology.

Signed-off-by: Timur Kristóf <timur.kristof@gmail.com>
Reviewed-by: Jason Ekstrand <jason.ekstrand@collabora.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/15103>

diff --git a/src/amd/vulkan/radv_shader.c b/src/amd/vulkan/radv_shader.c
index 4a5cae3..b0d3145 100644 (file)
--- a/src/amd/vulkan/radv_shader.c
+++ b/src/amd/vulkan/radv_shader.c
@@ -723,7 +723,7 @@ radv_shader_compile_to_nir(struct radv_device *device, struct vk_shader_module *
       /* Mesh shaders run as NGG which can implement local_invocation_index from
        * the wave ID in merged_wave_info, but they don't have local_invocation_ids.
        */
-      .lower_cs_local_id_from_index = nir->info.stage == MESA_SHADER_MESH,
+      .lower_cs_local_id_to_index = nir->info.stage == MESA_SHADER_MESH,
       .lower_local_invocation_index = nir->info.stage == MESA_SHADER_COMPUTE &&
                                       ((nir->info.workgroup_size[0] == 1) +
                                        (nir->info.workgroup_size[1] == 1) +

diff --git a/src/asahi/compiler/agx_compile.h b/src/asahi/compiler/agx_compile.h
index f1da074..8df4576 100644 (file)
--- a/src/asahi/compiler/agx_compile.h
+++ b/src/asahi/compiler/agx_compile.h
@@ -216,7 +216,7 @@ static const nir_shader_compiler_options agx_nir_options = {
    .lower_insert_byte = true,
    .lower_insert_word = true,
    .lower_uniforms_to_ubo = true,
-   .lower_cs_local_index_from_id = true,
+   .lower_cs_local_index_to_id = true,
 
    .lower_doubles_options = nir_lower_dmod,
    .lower_int64_options = ~(nir_lower_iadd64 | nir_lower_imul_2x32_64),

diff --git a/src/broadcom/vulkan/v3dv_pipeline.c b/src/broadcom/vulkan/v3dv_pipeline.c
index 0f4d43d..31d5b6e 100644 (file)
--- a/src/broadcom/vulkan/v3dv_pipeline.c
+++ b/src/broadcom/vulkan/v3dv_pipeline.c
@@ -204,7 +204,7 @@ const nir_shader_compiler_options v3dv_nir_options = {
    .lower_bitfield_extract_to_shifts = true,
    .lower_bitfield_reverse = true,
    .lower_bit_count = true,
-   .lower_cs_local_id_from_index = true,
+   .lower_cs_local_id_to_index = true,
    .lower_ffract = true,
    .lower_fmod = true,
    .lower_pack_unorm_2x16 = true,

diff --git a/src/compiler/nir/nir.h b/src/compiler/nir/nir.h
index a48fbb2..a1904e7 100644 (file)
--- a/src/compiler/nir/nir.h
+++ b/src/compiler/nir/nir.h
@@ -3337,8 +3337,8 @@ typedef struct nir_shader_compiler_options {
     */
    bool optimize_sample_mask_in;
 
-   bool lower_cs_local_index_from_id;
-   bool lower_cs_local_id_from_index;
+   bool lower_cs_local_index_to_id;
+   bool lower_cs_local_id_to_index;
 
    /* Prevents lowering global_invocation_id to be in terms of workgroup_id */
    bool has_cs_global_id;
@@ -4763,7 +4763,7 @@ typedef struct nir_lower_compute_system_values_options {
    bool has_base_workgroup_id:1;
    bool shuffle_local_ids_for_quad_derivatives:1;
    bool lower_local_invocation_index:1;
-   bool lower_cs_local_id_from_index:1;
+   bool lower_cs_local_id_to_index:1;
 } nir_lower_compute_system_values_options;
 
 bool nir_lower_compute_system_values(nir_shader *shader,

diff --git a/src/compiler/nir/nir_lower_system_values.c b/src/compiler/nir/nir_lower_system_values.c
index 37657b0..a2da7e5 100644 (file)
--- a/src/compiler/nir/nir_lower_system_values.c
+++ b/src/compiler/nir/nir_lower_system_values.c
@@ -288,12 +288,12 @@ lower_compute_system_value_instr(nir_builder *b,
 
    switch (intrin->intrinsic) {
    case nir_intrinsic_load_local_invocation_id:
-      /* If lower_cs_local_id_from_index is true, then we derive the local
-       * index from the local id.
+      /* If lower_cs_local_id_to_index is true, then we replace
+       * local_invocation_id with a formula based on local_invocation_index.
        */
-      if (b->shader->options->lower_cs_local_id_from_index ||
-          (options && options->lower_cs_local_id_from_index)) {
-         /* We lower gl_LocalInvocationID from gl_LocalInvocationIndex based
+      if (b->shader->options->lower_cs_local_id_to_index ||
+          (options && options->lower_cs_local_id_to_index)) {
+         /* We lower gl_LocalInvocationID to gl_LocalInvocationIndex based
           * on this formula:
           *
           *    gl_LocalInvocationID.x =
@@ -422,10 +422,10 @@ lower_compute_system_value_instr(nir_builder *b,
       return NULL;
 
    case nir_intrinsic_load_local_invocation_index:
-      /* If lower_cs_local_index_from_id is true, then we derive the local
-       * index from the local id.
+      /* If lower_cs_local_index_to_id is true, then we replace
+       * local_invocation_index with a formula based on local_invocation_id.
        */
-      if (b->shader->options->lower_cs_local_index_from_id ||
+      if (b->shader->options->lower_cs_local_index_to_id ||
           (options && options->lower_local_invocation_index)) {
          /* From the GLSL man page for gl_LocalInvocationIndex:
           *

diff --git a/src/freedreno/ir3/ir3_compiler.c b/src/freedreno/ir3/ir3_compiler.c
index c23d85e..3c9e0db 100644 (file)
--- a/src/freedreno/ir3/ir3_compiler.c
+++ b/src/freedreno/ir3/ir3_compiler.c
@@ -113,7 +113,7 @@ static const nir_shader_compiler_options options = {
    .has_fsub = true,
    .has_isub = true,
    .lower_wpos_pntc = true,
-   .lower_cs_local_index_from_id = true,
+   .lower_cs_local_index_to_id = true,
 
    /* Only needed for the spirv_to_nir() pass done in ir3_cmdline.c
     * but that should be harmless for GL since 64b is not
@@ -172,7 +172,7 @@ static const nir_shader_compiler_options options_a6xx = {
    .max_unroll_iterations = 32,
    .force_indirect_unrolling = nir_var_all,
    .lower_wpos_pntc = true,
-   .lower_cs_local_index_from_id = true,
+   .lower_cs_local_index_to_id = true,
 
    /* Only needed for the spirv_to_nir() pass done in ir3_cmdline.c
     * but that should be harmless for GL since 64b is not

diff --git a/src/gallium/drivers/i915/i915_screen.c b/src/gallium/drivers/i915/i915_screen.c
index 5ef24df..cfe1096 100644 (file)
--- a/src/gallium/drivers/i915/i915_screen.c
+++ b/src/gallium/drivers/i915/i915_screen.c
@@ -160,7 +160,7 @@ static const struct nir_shader_compiler_options gallivm_nir_options = {
    .lower_ifind_msb = true,
    .max_unroll_iterations = 32,
    .use_interpolated_input_intrinsics = true,
-   .lower_cs_local_index_from_id = true,
+   .lower_cs_local_index_to_id = true,
    .lower_uniforms_to_ubo = true,
    .lower_vector_cmp = true,
    .lower_device_index_to_zero = true,

diff --git a/src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp b/src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp
index 45cf808..87f6b67 100644 (file)
--- a/src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp
+++ b/src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp
@@ -3288,8 +3288,8 @@ nvir_nir_shader_compiler_options(int chipset)
    op.lower_base_vertex = false;
    op.lower_helper_invocation = false;
    op.optimize_sample_mask_in = false;
-   op.lower_cs_local_index_from_id = true;
-   op.lower_cs_local_id_from_index = false;
+   op.lower_cs_local_index_to_id = true;
+   op.lower_cs_local_id_to_index = false;
    op.lower_device_index_to_zero = false; // TODO
    op.lower_wpos_pntc = false; // TODO
    op.lower_hadd = true; // TODO

diff --git a/src/gallium/drivers/v3d/v3d_screen.c b/src/gallium/drivers/v3d/v3d_screen.c
index 6aa4568..955e198 100644 (file)
--- a/src/gallium/drivers/v3d/v3d_screen.c
+++ b/src/gallium/drivers/v3d/v3d_screen.c
@@ -695,7 +695,7 @@ static const nir_shader_compiler_options v3d_nir_options = {
         .lower_bitfield_extract_to_shifts = true,
         .lower_bitfield_reverse = true,
         .lower_bit_count = true,
-        .lower_cs_local_id_from_index = true,
+        .lower_cs_local_id_to_index = true,
         .lower_ffract = true,
         .lower_fmod = true,
         .lower_pack_unorm_2x16 = true,

diff --git a/src/panfrost/bifrost/bifrost_compile.h b/src/panfrost/bifrost/bifrost_compile.h
index 9b320d7..446ef24 100644 (file)
--- a/src/panfrost/bifrost/bifrost_compile.h
+++ b/src/panfrost/bifrost/bifrost_compile.h
@@ -91,7 +91,7 @@ static const nir_shader_compiler_options bifrost_nir_options = {
 
         .has_cs_global_id = true,
         .vertex_id_zero_based = true,
-        .lower_cs_local_index_from_id = true,
+        .lower_cs_local_index_to_id = true,
         .max_unroll_iterations = 32,
         .force_indirect_unrolling = (nir_var_shader_in | nir_var_shader_out | nir_var_function_temp),
 };

diff --git a/src/panfrost/midgard/midgard_compile.h b/src/panfrost/midgard/midgard_compile.h
index b24c63b..f9c44f9 100644 (file)
--- a/src/panfrost/midgard/midgard_compile.h
+++ b/src/panfrost/midgard/midgard_compile.h
@@ -94,7 +94,7 @@ static const nir_shader_compiler_options midgard_nir_options = {
 
         .vertex_id_zero_based = true,
         .has_cs_global_id = true,
-        .lower_cs_local_index_from_id = true,
+        .lower_cs_local_index_to_id = true,
         .max_unroll_iterations = 32,
         .force_indirect_unrolling = (nir_var_shader_in | nir_var_shader_out | nir_var_function_temp),
 };