nir: Rename WORK_GROUP (and similar) to WORKGROUP

Be consistent with other usages in Vulkan and SPIR-V, and the recently
added workgroup_size field.

Acked-by: Emma Anholt <emma@anholt.net>
Acked-by: Alyssa Rosenzweig <alyssa.rosenzweig@collabora.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Acked-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11190>

13 files changed, 61 insertions, 61 deletions

diff --git a/src/compiler/glsl/builtin_variables.cpp b/src/compiler/glsl/builtin_variables.cpp
index a314e32d029..3a8ec615c13 100644
--- a/src/compiler/glsl/builtin_variables.cpp
+++ b/src/compiler/glsl/builtin_variables.cpp
@@ -1452,8 +1452,8 @@ builtin_variable_generator::generate_cs_special_vars()
 {
    add_system_value(SYSTEM_VALUE_LOCAL_INVOCATION_ID, uvec3_t,
                     "gl_LocalInvocationID");
-   add_system_value(SYSTEM_VALUE_WORK_GROUP_ID, uvec3_t, "gl_WorkGroupID");
-   add_system_value(SYSTEM_VALUE_NUM_WORK_GROUPS, uvec3_t, "gl_NumWorkGroups");
+   add_system_value(SYSTEM_VALUE_WORKGROUP_ID, uvec3_t, "gl_WorkGroupID");
+   add_system_value(SYSTEM_VALUE_NUM_WORKGROUPS, uvec3_t, "gl_NumWorkGroups");
 
    if (state->ARB_compute_variable_group_size_enable) {
       add_system_value(SYSTEM_VALUE_WORKGROUP_SIZE,
diff --git a/src/compiler/glsl/lower_cs_derived.cpp b/src/compiler/glsl/lower_cs_derived.cpp
index 85a57232b61..99a0028fb6a 100644
--- a/src/compiler/glsl/lower_cs_derived.cpp
+++ b/src/compiler/glsl/lower_cs_derived.cpp
@@ -129,7 +129,7 @@ lower_cs_derived_visitor::find_sysvals()
 
    if (!gl_WorkGroupID)
       gl_WorkGroupID = add_system_value(
-            SYSTEM_VALUE_WORK_GROUP_ID, glsl_type::uvec3_type, "gl_WorkGroupID");
+            SYSTEM_VALUE_WORKGROUP_ID, glsl_type::uvec3_type, "gl_WorkGroupID");
    if (!gl_LocalInvocationID)
       gl_LocalInvocationID = add_system_value(
             SYSTEM_VALUE_LOCAL_INVOCATION_ID, glsl_type::uvec3_type,
diff --git a/src/compiler/nir/nir.c b/src/compiler/nir/nir.c
index 2a78308d809..2e45709dfcf 100644
--- a/src/compiler/nir/nir.c
+++ b/src/compiler/nir/nir.c
@@ -1975,10 +1975,10 @@ nir_intrinsic_from_system_value(gl_system_value val)
       return nir_intrinsic_load_local_invocation_id;
    case SYSTEM_VALUE_LOCAL_INVOCATION_INDEX:
       return nir_intrinsic_load_local_invocation_index;
-   case SYSTEM_VALUE_WORK_GROUP_ID:
-      return nir_intrinsic_load_work_group_id;
-   case SYSTEM_VALUE_NUM_WORK_GROUPS:
-      return nir_intrinsic_load_num_work_groups;
+   case SYSTEM_VALUE_WORKGROUP_ID:
+      return nir_intrinsic_load_workgroup_id;
+   case SYSTEM_VALUE_NUM_WORKGROUPS:
+      return nir_intrinsic_load_num_workgroups;
    case SYSTEM_VALUE_PRIMITIVE_ID:
       return nir_intrinsic_load_primitive_id;
    case SYSTEM_VALUE_TESS_COORD:
@@ -2106,10 +2106,10 @@ nir_system_value_from_intrinsic(nir_intrinsic_op intrin)
       return SYSTEM_VALUE_LOCAL_INVOCATION_ID;
    case nir_intrinsic_load_local_invocation_index:
       return SYSTEM_VALUE_LOCAL_INVOCATION_INDEX;
-   case nir_intrinsic_load_num_work_groups:
-      return SYSTEM_VALUE_NUM_WORK_GROUPS;
-   case nir_intrinsic_load_work_group_id:
-      return SYSTEM_VALUE_WORK_GROUP_ID;
+   case nir_intrinsic_load_num_workgroups:
+      return SYSTEM_VALUE_NUM_WORKGROUPS;
+   case nir_intrinsic_load_workgroup_id:
+      return SYSTEM_VALUE_WORKGROUP_ID;
    case nir_intrinsic_load_primitive_id:
       return SYSTEM_VALUE_PRIMITIVE_ID;
    case nir_intrinsic_load_tess_coord:
diff --git a/src/compiler/nir/nir.h b/src/compiler/nir/nir.h
index cedbc5bb32b..7981031b54a 100644
--- a/src/compiler/nir/nir.h
+++ b/src/compiler/nir/nir.h
@@ -3297,7 +3297,7 @@ typedef struct nir_shader_compiler_options {
    bool lower_cs_local_index_from_id;
    bool lower_cs_local_id_from_index;
 
-   /* Prevents lowering global_invocation_id to be in terms of work_group_id */
+   /* Prevents lowering global_invocation_id to be in terms of workgroup_id */
    bool has_cs_global_id;
 
    bool lower_device_index_to_zero;
@@ -4674,7 +4674,7 @@ bool nir_lower_system_values(nir_shader *shader);
 
 typedef struct nir_lower_compute_system_values_options {
    bool has_base_global_invocation_id:1;
-   bool has_base_work_group_id:1;
+   bool has_base_workgroup_id:1;
    bool shuffle_local_ids_for_quad_derivatives:1;
    bool lower_local_invocation_index:1;
 } nir_lower_compute_system_values_options;
@@ -5256,9 +5256,9 @@ nir_variable_is_in_block(const nir_variable *var)
 typedef struct nir_unsigned_upper_bound_config {
    unsigned min_subgroup_size;
    unsigned max_subgroup_size;
-   unsigned max_work_group_invocations;
-   unsigned max_work_group_count[3];
-   unsigned max_work_group_size[3];
+   unsigned max_workgroup_invocations;
+   unsigned max_workgroup_count[3];
+   unsigned max_workgroup_size[3];
 
    uint32_t vertex_attrib_max[32];
 } nir_unsigned_upper_bound_config;
diff --git a/src/compiler/nir/nir_divergence_analysis.c b/src/compiler/nir/nir_divergence_analysis.c
index 05d525fe803..5e74846232e 100644
--- a/src/compiler/nir/nir_divergence_analysis.c
+++ b/src/compiler/nir/nir_divergence_analysis.c
@@ -101,7 +101,7 @@ visit_intrinsic(nir_shader *shader, nir_intrinsic_instr *instr)
    case nir_intrinsic_vote_ieq:
    case nir_intrinsic_load_push_constant:
    case nir_intrinsic_load_work_dim:
-   case nir_intrinsic_load_num_work_groups:
+   case nir_intrinsic_load_num_workgroups:
    case nir_intrinsic_load_workgroup_size:
    case nir_intrinsic_load_subgroup_id:
    case nir_intrinsic_load_num_subgroups:
@@ -236,7 +236,7 @@ visit_intrinsic(nir_shader *shader, nir_intrinsic_instr *instr)
          assert(stage == MESA_SHADER_TESS_CTRL);
       break;
 
-   case nir_intrinsic_load_work_group_id:
+   case nir_intrinsic_load_workgroup_id:
       assert(stage == MESA_SHADER_COMPUTE);
       is_divergent |= (options & nir_divergence_multiple_workgroup_per_compute_subgroup);
       break;
diff --git a/src/compiler/nir/nir_gather_info.c b/src/compiler/nir/nir_gather_info.c
index bdf84a3a60d..5067a7eb861 100644
--- a/src/compiler/nir/nir_gather_info.c
+++ b/src/compiler/nir/nir_gather_info.c
@@ -623,8 +623,8 @@ gather_intrinsic_info(nir_intrinsic_instr *instr, nir_shader *shader,
    case nir_intrinsic_load_global_invocation_id:
    case nir_intrinsic_load_base_global_invocation_id:
    case nir_intrinsic_load_global_invocation_index:
-   case nir_intrinsic_load_work_group_id:
-   case nir_intrinsic_load_num_work_groups:
+   case nir_intrinsic_load_workgroup_id:
+   case nir_intrinsic_load_num_workgroups:
    case nir_intrinsic_load_workgroup_size:
    case nir_intrinsic_load_work_dim:
    case nir_intrinsic_load_user_data_amd:
diff --git a/src/compiler/nir/nir_intrinsics.py b/src/compiler/nir/nir_intrinsics.py
index 89f8992ea2c..5cb2f60637a 100644
--- a/src/compiler/nir/nir_intrinsics.py
+++ b/src/compiler/nir/nir_intrinsics.py
@@ -685,11 +685,11 @@ system_value("local_invocation_id", 3)
 system_value("local_invocation_index", 1)
 # zero_base indicates it starts from 0 for the current dispatch
 # non-zero_base indicates the base is included
-system_value("work_group_id", 3, bit_sizes=[32, 64])
-system_value("work_group_id_zero_base", 3)
-system_value("base_work_group_id", 3, bit_sizes=[32, 64])
+system_value("workgroup_id", 3, bit_sizes=[32, 64])
+system_value("workgroup_id_zero_base", 3)
+system_value("base_workgroup_id", 3, bit_sizes=[32, 64])
 system_value("user_clip_plane", 4, indices=[UCP_ID])
-system_value("num_work_groups", 3, bit_sizes=[32, 64])
+system_value("num_workgroups", 3, bit_sizes=[32, 64])
 system_value("helper_invocation", 1, bit_sizes=[1, 32])
 system_value("layer_id", 1)
 system_value("view_index", 1)
@@ -704,8 +704,8 @@ system_value("num_subgroups", 1)
 system_value("subgroup_id", 1)
 system_value("workgroup_size", 3)
 # note: the definition of global_invocation_id_zero_base is based on
-# (work_group_id * workgroup_size) + local_invocation_id.
-# it is *not* based on work_group_id_zero_base, meaning the work group
+# (workgroup_id * workgroup_size) + local_invocation_id.
+# it is *not* based on workgroup_id_zero_base, meaning the work group
 # base is already accounted for, and the global base is additive on top of that
 system_value("global_invocation_id", 3, bit_sizes=[32, 64])
 system_value("global_invocation_id_zero_base", 3, bit_sizes=[32, 64])
diff --git a/src/compiler/nir/nir_lower_system_values.c b/src/compiler/nir/nir_lower_system_values.c
index 5db6c2b6d40..732b8e80a6b 100644
--- a/src/compiler/nir/nir_lower_system_values.c
+++ b/src/compiler/nir/nir_lower_system_values.c
@@ -55,9 +55,9 @@ static nir_ssa_def*
 build_global_group_size(nir_builder *b, unsigned bit_size)
 {
    nir_ssa_def *group_size = nir_load_workgroup_size(b);
-   nir_ssa_def *num_work_groups = nir_load_num_work_groups(b, bit_size);
+   nir_ssa_def *num_workgroups = nir_load_num_workgroups(b, bit_size);
    return nir_imul(b, nir_u2u(b, group_size, bit_size),
-                      num_work_groups);
+                      num_workgroups);
 }
 
 static bool
@@ -443,10 +443,10 @@ lower_compute_system_value_instr(nir_builder *b,
       }
 
    case nir_intrinsic_load_global_invocation_id_zero_base: {
-      if ((options && options->has_base_work_group_id) ||
+      if ((options && options->has_base_workgroup_id) ||
           !b->shader->options->has_cs_global_id) {
          nir_ssa_def *group_size = nir_load_workgroup_size(b);
-         nir_ssa_def *group_id = nir_load_work_group_id(b, bit_size);
+         nir_ssa_def *group_id = nir_load_workgroup_id(b, bit_size);
          nir_ssa_def *local_id = nir_load_local_invocation_id(b);
 
          return nir_iadd(b, nir_imul(b, group_id,
@@ -461,7 +461,7 @@ lower_compute_system_value_instr(nir_builder *b,
       if (options && options->has_base_global_invocation_id)
          return nir_iadd(b, nir_load_global_invocation_id_zero_base(b, bit_size),
                             nir_load_base_global_invocation_id(b, bit_size));
-      else if ((options && options->has_base_work_group_id) ||
+      else if ((options && options->has_base_workgroup_id) ||
                !b->shader->options->has_cs_global_id)
          return nir_load_global_invocation_id_zero_base(b, bit_size);
       else
@@ -485,10 +485,10 @@ lower_compute_system_value_instr(nir_builder *b,
       return index;
    }
 
-   case nir_intrinsic_load_work_group_id: {
-      if (options && options->has_base_work_group_id)
-         return nir_iadd(b, nir_u2u(b, nir_load_work_group_id_zero_base(b), bit_size),
-                            nir_load_base_work_group_id(b, bit_size));
+   case nir_intrinsic_load_workgroup_id: {
+      if (options && options->has_base_workgroup_id)
+         return nir_iadd(b, nir_u2u(b, nir_load_workgroup_id_zero_base(b), bit_size),
+                            nir_load_base_workgroup_id(b, bit_size));
       else
          return NULL;
    }
diff --git a/src/compiler/nir/nir_opt_peephole_select.c b/src/compiler/nir/nir_opt_peephole_select.c
index e3cb21be2d5..62530ddd793 100644
--- a/src/compiler/nir/nir_opt_peephole_select.c
+++ b/src/compiler/nir/nir_opt_peephole_select.c
@@ -106,8 +106,8 @@ block_check_for_allowed_instrs(nir_block *block, unsigned *count,
          case nir_intrinsic_load_base_instance:
          case nir_intrinsic_load_instance_id:
          case nir_intrinsic_load_draw_id:
-         case nir_intrinsic_load_num_work_groups:
-         case nir_intrinsic_load_work_group_id:
+         case nir_intrinsic_load_num_workgroups:
+         case nir_intrinsic_load_workgroup_id:
          case nir_intrinsic_load_local_invocation_id:
          case nir_intrinsic_load_local_invocation_index:
          case nir_intrinsic_load_subgroup_id:
diff --git a/src/compiler/nir/nir_range_analysis.c b/src/compiler/nir/nir_range_analysis.c
index 608e1c0dd28..18ac161d5e0 100644
--- a/src/compiler/nir/nir_range_analysis.c
+++ b/src/compiler/nir/nir_range_analysis.c
@@ -1256,9 +1256,9 @@ lookup_input(nir_shader *shader, unsigned driver_location)
 static const nir_unsigned_upper_bound_config default_ub_config = {
    .min_subgroup_size = 1u,
    .max_subgroup_size = UINT16_MAX,
-   .max_work_group_invocations = UINT16_MAX,
-   .max_work_group_count = {UINT16_MAX, UINT16_MAX, UINT16_MAX},
-   .max_work_group_size = {UINT16_MAX, UINT16_MAX, UINT16_MAX},
+   .max_workgroup_invocations = UINT16_MAX,
+   .max_workgroup_count = {UINT16_MAX, UINT16_MAX, UINT16_MAX},
+   .max_workgroup_size = {UINT16_MAX, UINT16_MAX, UINT16_MAX},
    .vertex_attrib_max = {
       UINT32_MAX, UINT32_MAX, UINT32_MAX, UINT32_MAX, UINT32_MAX, UINT32_MAX, UINT32_MAX, UINT32_MAX,
       UINT32_MAX, UINT32_MAX, UINT32_MAX, UINT32_MAX, UINT32_MAX, UINT32_MAX, UINT32_MAX, UINT32_MAX,
@@ -1294,7 +1294,7 @@ nir_unsigned_upper_bound(nir_shader *shader, struct hash_table *range_ht,
       case nir_intrinsic_load_local_invocation_index:
          if (shader->info.stage != MESA_SHADER_COMPUTE ||
              shader->info.cs.workgroup_size_variable) {
-            res = config->max_work_group_invocations - 1;
+            res = config->max_workgroup_invocations - 1;
          } else {
             res = (shader->info.cs.workgroup_size[0] *
                    shader->info.cs.workgroup_size[1] *
@@ -1303,23 +1303,23 @@ nir_unsigned_upper_bound(nir_shader *shader, struct hash_table *range_ht,
          break;
       case nir_intrinsic_load_local_invocation_id:
          if (shader->info.cs.workgroup_size_variable)
-            res = config->max_work_group_size[scalar.comp] - 1u;
+            res = config->max_workgroup_size[scalar.comp] - 1u;
          else
             res = shader->info.cs.workgroup_size[scalar.comp] - 1u;
          break;
-      case nir_intrinsic_load_work_group_id:
-         res = config->max_work_group_count[scalar.comp] - 1u;
+      case nir_intrinsic_load_workgroup_id:
+         res = config->max_workgroup_count[scalar.comp] - 1u;
          break;
-      case nir_intrinsic_load_num_work_groups:
-         res = config->max_work_group_count[scalar.comp];
+      case nir_intrinsic_load_num_workgroups:
+         res = config->max_workgroup_count[scalar.comp];
          break;
       case nir_intrinsic_load_global_invocation_id:
          if (shader->info.cs.workgroup_size_variable) {
-            res = mul_clamp(config->max_work_group_size[scalar.comp],
-                            config->max_work_group_count[scalar.comp]) - 1u;
+            res = mul_clamp(config->max_workgroup_size[scalar.comp],
+                            config->max_workgroup_count[scalar.comp]) - 1u;
          } else {
             res = (shader->info.cs.workgroup_size[scalar.comp] *
-                   config->max_work_group_count[scalar.comp]) - 1u;
+                   config->max_workgroup_count[scalar.comp]) - 1u;
          }
          break;
       case nir_intrinsic_load_invocation_id:
@@ -1338,13 +1338,13 @@ nir_unsigned_upper_bound(nir_shader *shader, struct hash_table *range_ht,
          break;
       case nir_intrinsic_load_subgroup_id:
       case nir_intrinsic_load_num_subgroups: {
-         uint32_t work_group_size = config->max_work_group_invocations;
+         uint32_t workgroup_size = config->max_workgroup_invocations;
          if (shader->info.stage == MESA_SHADER_COMPUTE && !shader->info.cs.workgroup_size_variable) {
-            work_group_size = shader->info.cs.workgroup_size[0] *
-                              shader->info.cs.workgroup_size[1] *
-                              shader->info.cs.workgroup_size[2];
+            workgroup_size = shader->info.cs.workgroup_size[0] *
+                             shader->info.cs.workgroup_size[1] *
+                             shader->info.cs.workgroup_size[2];
          }
-         res = DIV_ROUND_UP(work_group_size, config->min_subgroup_size);
+         res = DIV_ROUND_UP(workgroup_size, config->min_subgroup_size);
          if (intrin->intrinsic == nir_intrinsic_load_subgroup_id)
             res--;
          break;
@@ -1391,7 +1391,7 @@ nir_unsigned_upper_bound(nir_shader *shader, struct hash_table *range_ht,
       case nir_intrinsic_load_tess_rel_patch_id_amd:
       case nir_intrinsic_load_tcs_num_patches_amd:
          /* Very generous maximum: TCS/TES executed by largest possible workgroup */
-         res = config->max_work_group_invocations / MAX2(shader->info.tess.tcs_vertices_out, 1u);
+         res = config->max_workgroup_invocations / MAX2(shader->info.tess.tcs_vertices_out, 1u);
          break;
       default:
          break;
diff --git a/src/compiler/shader_enums.c b/src/compiler/shader_enums.c
index b1eae07c86e..c8529cd857b 100644
--- a/src/compiler/shader_enums.c
+++ b/src/compiler/shader_enums.c
@@ -275,8 +275,8 @@ gl_system_value_name(gl_system_value sysval)
      ENUM(SYSTEM_VALUE_GLOBAL_INVOCATION_ID),
      ENUM(SYSTEM_VALUE_BASE_GLOBAL_INVOCATION_ID),
      ENUM(SYSTEM_VALUE_GLOBAL_INVOCATION_INDEX),
-     ENUM(SYSTEM_VALUE_WORK_GROUP_ID),
-     ENUM(SYSTEM_VALUE_NUM_WORK_GROUPS),
+     ENUM(SYSTEM_VALUE_WORKGROUP_ID),
+     ENUM(SYSTEM_VALUE_NUM_WORKGROUPS),
      ENUM(SYSTEM_VALUE_WORKGROUP_SIZE),
      ENUM(SYSTEM_VALUE_GLOBAL_GROUP_SIZE),
      ENUM(SYSTEM_VALUE_USER_DATA_AMD),
diff --git a/src/compiler/shader_enums.h b/src/compiler/shader_enums.h
index 541742517e1..9350459a100 100644
--- a/src/compiler/shader_enums.h
+++ b/src/compiler/shader_enums.h
@@ -710,8 +710,8 @@ typedef enum
    SYSTEM_VALUE_GLOBAL_INVOCATION_ID,
    SYSTEM_VALUE_BASE_GLOBAL_INVOCATION_ID,
    SYSTEM_VALUE_GLOBAL_INVOCATION_INDEX,
-   SYSTEM_VALUE_WORK_GROUP_ID,
-   SYSTEM_VALUE_NUM_WORK_GROUPS,
+   SYSTEM_VALUE_WORKGROUP_ID,
+   SYSTEM_VALUE_NUM_WORKGROUPS,
    SYSTEM_VALUE_WORKGROUP_SIZE,
    SYSTEM_VALUE_GLOBAL_GROUP_SIZE,
    SYSTEM_VALUE_WORK_DIM,
diff --git a/src/compiler/spirv/vtn_variables.c b/src/compiler/spirv/vtn_variables.c
index c3b5fc7e4ff..a93d3d52067 100644
--- a/src/compiler/spirv/vtn_variables.c
+++ b/src/compiler/spirv/vtn_variables.c
@@ -914,7 +914,7 @@ vtn_get_builtin_location(struct vtn_builder *b,
       set_mode_system_value(b, mode);
       break;
    case SpvBuiltInNumWorkgroups:
-      *location = SYSTEM_VALUE_NUM_WORK_GROUPS;
+      *location = SYSTEM_VALUE_NUM_WORKGROUPS;
       set_mode_system_value(b, mode);
       break;
    case SpvBuiltInWorkgroupSize:
@@ -923,7 +923,7 @@ vtn_get_builtin_location(struct vtn_builder *b,
       set_mode_system_value(b, mode);
       break;
    case SpvBuiltInWorkgroupId:
-      *location = SYSTEM_VALUE_WORK_GROUP_ID;
+      *location = SYSTEM_VALUE_WORKGROUP_ID;
       set_mode_system_value(b, mode);
       break;
    case SpvBuiltInLocalInvocationId: