diff --git a/src/amd/compiler/aco_instruction_selection.cpp b/src/amd/compiler/aco_instruction_selection.cpp
index f4671a45c1f..9e1a26c79cc 100644
--- a/src/amd/compiler/aco_instruction_selection.cpp
+++ b/src/amd/compiler/aco_instruction_selection.cpp
@@ -7818,7 +7818,7 @@ void visit_intrinsic(isel_context *ctx, nir_intrinsic_instr *instr)
    case nir_intrinsic_scoped_barrier:
       emit_scoped_barrier(ctx, instr);
       break;
-   case nir_intrinsic_load_num_work_groups: {
+   case nir_intrinsic_load_num_workgroups: {
       Temp dst = get_ssa_temp(ctx, &instr->dest.ssa);
       bld.copy(Definition(dst), Operand(get_arg(ctx, ctx->args->ac.num_work_groups)));
       emit_split_vector(ctx, dst, 3);
@@ -7830,7 +7830,7 @@ void visit_intrinsic(isel_context *ctx, nir_intrinsic_instr *instr)
       emit_split_vector(ctx, dst, 3);
       break;
    }
-   case nir_intrinsic_load_work_group_id: {
+   case nir_intrinsic_load_workgroup_id: {
       Temp dst = get_ssa_temp(ctx, &instr->dest.ssa);
       struct ac_arg *args = ctx->args->ac.workgroup_ids;
       bld.pseudo(aco_opcode::p_create_vector, Definition(dst),
diff --git a/src/amd/compiler/aco_instruction_selection_setup.cpp b/src/amd/compiler/aco_instruction_selection_setup.cpp
index 36d5640f21c..f60937fc9c0 100644
--- a/src/amd/compiler/aco_instruction_selection_setup.cpp
+++ b/src/amd/compiler/aco_instruction_selection_setup.cpp
@@ -530,13 +530,13 @@ void init_context(isel_context *ctx, nir_shader *shader)
       ctx->ub_config.min_subgroup_size = ctx->options->key.cs.subgroup_size;
       ctx->ub_config.max_subgroup_size = ctx->options->key.cs.subgroup_size;
    }
-   ctx->ub_config.max_work_group_invocations = 2048;
-   ctx->ub_config.max_work_group_count[0] = 65535;
-   ctx->ub_config.max_work_group_count[1] = 65535;
-   ctx->ub_config.max_work_group_count[2] = 65535;
-   ctx->ub_config.max_work_group_size[0] = 2048;
-   ctx->ub_config.max_work_group_size[1] = 2048;
-   ctx->ub_config.max_work_group_size[2] = 2048;
+   ctx->ub_config.max_workgroup_invocations = 2048;
+   ctx->ub_config.max_workgroup_count[0] = 65535;
+   ctx->ub_config.max_workgroup_count[1] = 65535;
+   ctx->ub_config.max_workgroup_count[2] = 65535;
+   ctx->ub_config.max_workgroup_size[0] = 2048;
+   ctx->ub_config.max_workgroup_size[1] = 2048;
+   ctx->ub_config.max_workgroup_size[2] = 2048;
    for (unsigned i = 0; i < MAX_VERTEX_ATTRIBS; i++) {
       unsigned attrib_format = ctx->options->key.vs.vertex_attribute_formats[i];
       unsigned dfmt = attrib_format & 0xf;
@@ -719,8 +719,8 @@ void init_context(isel_context *ctx, nir_shader *shader)
                RegType type = RegType::sgpr;
                switch(intrinsic->intrinsic) {
                   case nir_intrinsic_load_push_constant:
-                  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_subgroup_id:
                   case nir_intrinsic_load_num_subgroups:
                   case nir_intrinsic_load_first_vertex:
diff --git a/src/amd/llvm/ac_nir_to_llvm.c b/src/amd/llvm/ac_nir_to_llvm.c
index 5fabf94757d..f2173376cce 100644
--- a/src/amd/llvm/ac_nir_to_llvm.c
+++ b/src/amd/llvm/ac_nir_to_llvm.c
@@ -3400,7 +3400,7 @@ static void visit_intrinsic(struct ac_nir_context *ctx, nir_intrinsic_instr *ins
    case nir_intrinsic_load_subgroup_invocation:
       result = ac_get_thread_id(&ctx->ac);
       break;
-   case nir_intrinsic_load_work_group_id: {
+   case nir_intrinsic_load_workgroup_id: {
       LLVMValueRef values[3];
 
       for (int i = 0; i < 3; i++) {
@@ -3516,7 +3516,7 @@ static void visit_intrinsic(struct ac_nir_context *ctx, nir_intrinsic_instr *ins
    case nir_intrinsic_load_instance_id:
       result = ctx->abi->instance_id;
       break;
-   case nir_intrinsic_load_num_work_groups:
+   case nir_intrinsic_load_num_workgroups:
       result = ac_get_arg(&ctx->ac, ctx->args->num_work_groups);
       break;
    case nir_intrinsic_load_local_invocation_index:
diff --git a/src/amd/vulkan/radv_meta_buffer.c b/src/amd/vulkan/radv_meta_buffer.c
index 8955438d1af..5f6e0afb829 100644
--- a/src/amd/vulkan/radv_meta_buffer.c
+++ b/src/amd/vulkan/radv_meta_buffer.c
@@ -13,7 +13,7 @@ build_buffer_fill_shader(struct radv_device *dev)
    b.shader->info.cs.workgroup_size[2] = 1;
 
    nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
-   nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
+   nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32);
    nir_ssa_def *block_size =
       nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
                     b.shader->info.cs.workgroup_size[2], 0);
@@ -43,7 +43,7 @@ build_buffer_copy_shader(struct radv_device *dev)
    b.shader->info.cs.workgroup_size[2] = 1;
 
    nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
-   nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
+   nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32);
    nir_ssa_def *block_size =
       nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
                     b.shader->info.cs.workgroup_size[2], 0);
diff --git a/src/amd/vulkan/radv_meta_bufimage.c b/src/amd/vulkan/radv_meta_bufimage.c
index ba272ed15a8..d472930932a 100644
--- a/src/amd/vulkan/radv_meta_bufimage.c
+++ b/src/amd/vulkan/radv_meta_bufimage.c
@@ -52,7 +52,7 @@ build_nir_itob_compute_shader(struct radv_device *dev, bool is_3d)
    output_img->data.binding = 1;
 
    nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
-   nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
+   nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32);
    nir_ssa_def *block_size =
       nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
                     b.shader->info.cs.workgroup_size[2], 0);
@@ -239,7 +239,7 @@ build_nir_btoi_compute_shader(struct radv_device *dev, bool is_3d)
    output_img->data.binding = 1;
 
    nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
-   nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
+   nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32);
    nir_ssa_def *block_size =
       nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
                     b.shader->info.cs.workgroup_size[2], 0);
@@ -421,7 +421,7 @@ build_nir_btoi_r32g32b32_compute_shader(struct radv_device *dev)
    output_img->data.binding = 1;
 
    nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
-   nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
+   nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32);
    nir_ssa_def *block_size =
       nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
                     b.shader->info.cs.workgroup_size[2], 0);
@@ -583,7 +583,7 @@ build_nir_itoi_compute_shader(struct radv_device *dev, bool is_3d, int samples)
    output_img->data.binding = 1;
 
    nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
-   nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
+   nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32);
    nir_ssa_def *block_size =
       nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
                     b.shader->info.cs.workgroup_size[2], 0);
@@ -785,7 +785,7 @@ build_nir_itoi_r32g32b32_compute_shader(struct radv_device *dev)
    output_img->data.binding = 1;
 
    nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
-   nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
+   nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32);
    nir_ssa_def *block_size =
       nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
                     b.shader->info.cs.workgroup_size[2], 0);
@@ -951,7 +951,7 @@ build_nir_cleari_compute_shader(struct radv_device *dev, bool is_3d, int samples
    output_img->data.binding = 0;
 
    nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
-   nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
+   nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32);
    nir_ssa_def *block_size =
       nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
                     b.shader->info.cs.workgroup_size[2], 0);
@@ -1117,7 +1117,7 @@ build_nir_cleari_r32g32b32_compute_shader(struct radv_device *dev)
    output_img->data.binding = 0;
 
    nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
-   nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
+   nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32);
    nir_ssa_def *block_size =
       nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
                     b.shader->info.cs.workgroup_size[2], 0);
diff --git a/src/amd/vulkan/radv_meta_clear.c b/src/amd/vulkan/radv_meta_clear.c
index 865aa35490d..44016c99e17 100644
--- a/src/amd/vulkan/radv_meta_clear.c
+++ b/src/amd/vulkan/radv_meta_clear.c
@@ -1007,7 +1007,7 @@ build_clear_htile_mask_shader()
    b.shader->info.cs.workgroup_size[2] = 1;
 
    nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
-   nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
+   nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32);
    nir_ssa_def *block_size =
       nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
                     b.shader->info.cs.workgroup_size[2], 0);
diff --git a/src/amd/vulkan/radv_meta_copy_vrs_htile.c b/src/amd/vulkan/radv_meta_copy_vrs_htile.c
index 7238fe25176..f8f15f89bf1 100644
--- a/src/amd/vulkan/radv_meta_copy_vrs_htile.c
+++ b/src/amd/vulkan/radv_meta_copy_vrs_htile.c
@@ -50,7 +50,7 @@ build_copy_vrs_htile_shader(struct radv_device *device, struct radeon_surf *surf
    b.shader->info.cs.workgroup_size[2] = 1;
 
    nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
-   nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
+   nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32);
    nir_ssa_def *block_size =
       nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
                     b.shader->info.cs.workgroup_size[2], 0);
diff --git a/src/amd/vulkan/radv_meta_dcc_retile.c b/src/amd/vulkan/radv_meta_dcc_retile.c
index 87c6cf60f89..8dba826e13f 100644
--- a/src/amd/vulkan/radv_meta_dcc_retile.c
+++ b/src/amd/vulkan/radv_meta_dcc_retile.c
@@ -33,7 +33,7 @@ get_global_ids(nir_builder *b, unsigned num_components)
    unsigned mask = BITFIELD_MASK(num_components);
 
    nir_ssa_def *local_ids = nir_channels(b, nir_load_local_invocation_id(b), mask);
-   nir_ssa_def *block_ids = nir_channels(b, nir_load_work_group_id(b, 32), mask);
+   nir_ssa_def *block_ids = nir_channels(b, nir_load_workgroup_id(b, 32), mask);
    nir_ssa_def *block_size = nir_channels(
       b,
       nir_imm_ivec4(b, b->shader->info.cs.workgroup_size[0], b->shader->info.cs.workgroup_size[1],
diff --git a/src/amd/vulkan/radv_meta_fast_clear.c b/src/amd/vulkan/radv_meta_fast_clear.c
index 27378de2da5..f7486ae6c67 100644
--- a/src/amd/vulkan/radv_meta_fast_clear.c
+++ b/src/amd/vulkan/radv_meta_fast_clear.c
@@ -49,7 +49,7 @@ build_dcc_decompress_compute_shader(struct radv_device *dev)
    output_img->data.binding = 1;
 
    nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
-   nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
+   nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32);
    nir_ssa_def *block_size =
       nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
                     b.shader->info.cs.workgroup_size[2], 0);
diff --git a/src/amd/vulkan/radv_meta_fmask_expand.c b/src/amd/vulkan/radv_meta_fmask_expand.c
index 7355b88979a..5bf7f26e659 100644
--- a/src/amd/vulkan/radv_meta_fmask_expand.c
+++ b/src/amd/vulkan/radv_meta_fmask_expand.c
@@ -49,7 +49,7 @@ build_fmask_expand_compute_shader(struct radv_device *device, int samples)
    output_img->data.access = ACCESS_NON_READABLE;
 
    nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
-   nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
+   nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32);
    nir_ssa_def *block_size =
       nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
                     b.shader->info.cs.workgroup_size[2], 0);
diff --git a/src/amd/vulkan/radv_meta_resolve_cs.c b/src/amd/vulkan/radv_meta_resolve_cs.c
index 027346db7a7..01df68a58a8 100644
--- a/src/amd/vulkan/radv_meta_resolve_cs.c
+++ b/src/amd/vulkan/radv_meta_resolve_cs.c
@@ -79,7 +79,7 @@ build_resolve_compute_shader(struct radv_device *dev, bool is_integer, bool is_s
    output_img->data.descriptor_set = 0;
    output_img->data.binding = 1;
    nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
-   nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
+   nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32);
    nir_ssa_def *block_size =
       nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
                     b.shader->info.cs.workgroup_size[2], 0);
@@ -149,7 +149,7 @@ build_depth_stencil_resolve_compute_shader(struct radv_device *dev, int samples,
    output_img->data.descriptor_set = 0;
    output_img->data.binding = 1;
    nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
-   nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
+   nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32);
    nir_ssa_def *block_size =
       nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
                     b.shader->info.cs.workgroup_size[2], 0);
diff --git a/src/amd/vulkan/radv_query.c b/src/amd/vulkan/radv_query.c
index 500c5089e65..5efb7dfb739 100644
--- a/src/amd/vulkan/radv_query.c
+++ b/src/amd/vulkan/radv_query.c
@@ -149,7 +149,7 @@ build_occlusion_query_shader(struct radv_device *device)
    nir_ssa_def *src_buf = radv_meta_load_descriptor(&b, 0, 1);
 
    nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
-   nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
+   nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32);
    nir_ssa_def *block_size =
       nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
                     b.shader->info.cs.workgroup_size[2], 0);
@@ -290,7 +290,7 @@ build_pipeline_statistics_query_shader(struct radv_device *device)
    nir_ssa_def *src_buf = radv_meta_load_descriptor(&b, 0, 1);
 
    nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
-   nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
+   nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32);
    nir_ssa_def *block_size =
       nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
                     b.shader->info.cs.workgroup_size[2], 0);
@@ -441,7 +441,7 @@ build_tfb_query_shader(struct radv_device *device)
 
    /* Compute global ID. */
    nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
-   nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
+   nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32);
    nir_ssa_def *block_size =
       nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
                     b.shader->info.cs.workgroup_size[2], 0);
@@ -571,7 +571,7 @@ build_timestamp_query_shader(struct radv_device *device)
 
    /* Compute global ID. */
    nir_ssa_def *invoc_id = nir_load_local_invocation_id(&b);
-   nir_ssa_def *wg_id = nir_load_work_group_id(&b, 32);
+   nir_ssa_def *wg_id = nir_load_workgroup_id(&b, 32);
    nir_ssa_def *block_size =
       nir_imm_ivec4(&b, b.shader->info.cs.workgroup_size[0], b.shader->info.cs.workgroup_size[1],
                     b.shader->info.cs.workgroup_size[2], 0);
diff --git a/src/amd/vulkan/radv_shader_info.c b/src/amd/vulkan/radv_shader_info.c
index fbac6ab10b0..6c77e02e383 100644
--- a/src/amd/vulkan/radv_shader_info.c
+++ b/src/amd/vulkan/radv_shader_info.c
@@ -158,16 +158,16 @@ gather_intrinsic_info(const nir_shader *nir, const nir_intrinsic_instr *instr,
    case nir_intrinsic_load_instance_id:
       info->vs.needs_instance_id = true;
       break;
-   case nir_intrinsic_load_num_work_groups:
+   case nir_intrinsic_load_num_workgroups:
       info->cs.uses_grid_size = true;
       break;
    case nir_intrinsic_load_local_invocation_id:
-   case nir_intrinsic_load_work_group_id: {
+   case nir_intrinsic_load_workgroup_id: {
       unsigned mask = nir_ssa_def_components_read(&instr->dest.ssa);
       while (mask) {
          unsigned i = u_bit_scan(&mask);
 
-         if (instr->intrinsic == nir_intrinsic_load_work_group_id)
+         if (instr->intrinsic == nir_intrinsic_load_workgroup_id)
             info->cs.uses_block_id[i] = true;
          else
             info->cs.uses_thread_id[i] = true;
diff --git a/src/broadcom/compiler/nir_to_vir.c b/src/broadcom/compiler/nir_to_vir.c
index 17475b33063..0cf66f5d4cd 100644
--- a/src/broadcom/compiler/nir_to_vir.c
+++ b/src/broadcom/compiler/nir_to_vir.c
@@ -3026,7 +3026,7 @@ ntq_emit_intrinsic(struct v3d_compile *c, nir_intrinsic_instr *instr)
                 vir_emit_thrsw(c);
                 break;
 
-        case nir_intrinsic_load_num_work_groups:
+        case nir_intrinsic_load_num_workgroups:
                 for (int i = 0; i < 3; i++) {
                         ntq_store_dest(c, &instr->dest, i,
                                        vir_uniform(c, QUNIFORM_NUM_WORK_GROUPS,
@@ -3040,7 +3040,7 @@ ntq_emit_intrinsic(struct v3d_compile *c, nir_intrinsic_instr *instr)
                                        vir_uniform_ui(c, 32 - c->local_invocation_index_bits)));
                 break;
 
-        case nir_intrinsic_load_work_group_id: {
+        case nir_intrinsic_load_workgroup_id: {
                 struct qreg x = vir_AND(c, c->cs_payload[0],
                                          vir_uniform_ui(c, 0xffff));
 
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:
diff --git a/src/freedreno/computerator/a6xx.c b/src/freedreno/computerator/a6xx.c
index 180c67ea840..b17d7dcac38 100644
--- a/src/freedreno/computerator/a6xx.c
+++ b/src/freedreno/computerator/a6xx.c
@@ -156,7 +156,7 @@ cs_program_emit(struct fd_ringbuffer *ring, struct kernel *kernel)
    uint32_t local_invocation_id, work_group_id;
    local_invocation_id =
       ir3_find_sysval_regid(v, SYSTEM_VALUE_LOCAL_INVOCATION_ID);
-   work_group_id = ir3_find_sysval_regid(v, SYSTEM_VALUE_WORK_GROUP_ID);
+   work_group_id = ir3_find_sysval_regid(v, SYSTEM_VALUE_WORKGROUP_ID);
 
    OUT_PKT4(ring, REG_A6XX_HLSQ_CS_CNTL_0, 2);
    OUT_RING(ring, A6XX_HLSQ_CS_CNTL_0_WGIDCONSTID(work_group_id) |
diff --git a/src/freedreno/ir3/ir3_compiler_nir.c b/src/freedreno/ir3/ir3_compiler_nir.c
index 1ff3aa353fb..c88c30543d2 100644
--- a/src/freedreno/ir3/ir3_compiler_nir.c
+++ b/src/freedreno/ir3/ir3_compiler_nir.c
@@ -2023,21 +2023,21 @@ emit_intrinsic(struct ir3_context *ctx, nir_intrinsic_instr *intr)
 		}
 		ir3_split_dest(b, dst, ctx->local_invocation_id, 0, 3);
 		break;
-	case nir_intrinsic_load_work_group_id:
-	case nir_intrinsic_load_work_group_id_zero_base:
+	case nir_intrinsic_load_workgroup_id:
+	case nir_intrinsic_load_workgroup_id_zero_base:
 		if (!ctx->work_group_id) {
 			ctx->work_group_id =
-				create_sysval_input(ctx, SYSTEM_VALUE_WORK_GROUP_ID, 0x7);
+				create_sysval_input(ctx, SYSTEM_VALUE_WORKGROUP_ID, 0x7);
 			ctx->work_group_id->regs[0]->flags |= IR3_REG_SHARED;
 		}
 		ir3_split_dest(b, dst, ctx->work_group_id, 0, 3);
 		break;
-	case nir_intrinsic_load_base_work_group_id:
+	case nir_intrinsic_load_base_workgroup_id:
 		for (int i = 0; i < dest_components; i++) {
 			dst[i] = create_driver_param(ctx, IR3_DP_BASE_GROUP_X + i);
 		}
 		break;
-	case nir_intrinsic_load_num_work_groups:
+	case nir_intrinsic_load_num_workgroups:
 		for (int i = 0; i < dest_components; i++) {
 			dst[i] = create_driver_param(ctx, IR3_DP_NUM_WORK_GROUPS_X + i);
 		}
diff --git a/src/freedreno/ir3/ir3_nir.c b/src/freedreno/ir3/ir3_nir.c
index 0353bb33b00..a45274b82df 100644
--- a/src/freedreno/ir3/ir3_nir.c
+++ b/src/freedreno/ir3/ir3_nir.c
@@ -676,7 +676,7 @@ ir3_nir_scan_driver_consts(nir_shader *shader,
 					layout->num_driver_params =
 						MAX2(layout->num_driver_params, IR3_DP_UCP0_X + (idx + 1) * 4);
 					break;
-				case nir_intrinsic_load_num_work_groups:
+				case nir_intrinsic_load_num_workgroups:
 					layout->num_driver_params =
 						MAX2(layout->num_driver_params, IR3_DP_NUM_WORK_GROUPS_Z + 1);
 					break;
@@ -684,7 +684,7 @@ ir3_nir_scan_driver_consts(nir_shader *shader,
 					layout->num_driver_params =
 						MAX2(layout->num_driver_params, IR3_DP_LOCAL_GROUP_SIZE_Z + 1);
 					break;
-				case nir_intrinsic_load_base_work_group_id:
+				case nir_intrinsic_load_base_workgroup_id:
 					layout->num_driver_params =
 						MAX2(layout->num_driver_params, IR3_DP_BASE_GROUP_Z + 1);
 					break;
diff --git a/src/freedreno/ir3/ir3_parser.y b/src/freedreno/ir3/ir3_parser.y
index ae23c6a1da0..6609f216b6a 100644
--- a/src/freedreno/ir3/ir3_parser.y
+++ b/src/freedreno/ir3/ir3_parser.y
@@ -671,7 +671,7 @@ wgid_header:       T_A_WGID '(' T_REGISTER ')' {
                        assert(($3 & 0x1) == 0);  /* half-reg not allowed */
                        unsigned reg = $3 >> 1;
                        assert(reg >= regid(48, 0)); /* must be a high reg */
-                       add_sysval(reg, 0x7, SYSTEM_VALUE_WORK_GROUP_ID);
+                       add_sysval(reg, 0x7, SYSTEM_VALUE_WORKGROUP_ID);
 }
 
 numwg_header:      T_A_NUMWG '(' T_CONSTANT ')' {
diff --git a/src/freedreno/vulkan/tu_pipeline.c b/src/freedreno/vulkan/tu_pipeline.c
index 977c1c54440..1017cd638ef 100644
--- a/src/freedreno/vulkan/tu_pipeline.c
+++ b/src/freedreno/vulkan/tu_pipeline.c
@@ -578,7 +578,7 @@ tu6_emit_cs_config(struct tu_cs *cs, const struct tu_shader *shader,
    uint32_t local_invocation_id =
       ir3_find_sysval_regid(v, SYSTEM_VALUE_LOCAL_INVOCATION_ID);
    uint32_t work_group_id =
-      ir3_find_sysval_regid(v, SYSTEM_VALUE_WORK_GROUP_ID);
+      ir3_find_sysval_regid(v, SYSTEM_VALUE_WORKGROUP_ID);
 
    enum a6xx_threadsize thrsz = v->info.double_threadsize ? THREAD128 : THREAD64;
    tu_cs_emit_pkt4(cs, REG_A6XX_HLSQ_CS_CNTL_0, 2);
diff --git a/src/freedreno/vulkan/tu_shader.c b/src/freedreno/vulkan/tu_shader.c
index 0991ec35240..967bd4357f3 100644
--- a/src/freedreno/vulkan/tu_shader.c
+++ b/src/freedreno/vulkan/tu_shader.c
@@ -84,7 +84,7 @@ tu_spirv_to_nir(struct tu_device *dev,
    };
 
    const struct nir_lower_compute_system_values_options compute_sysval_options = {
-      .has_base_work_group_id = true,
+      .has_base_workgroup_id = true,
    };
 
    const nir_shader_compiler_options *nir_options =
diff --git a/src/gallium/auxiliary/gallivm/lp_bld_nir.c b/src/gallium/auxiliary/gallivm/lp_bld_nir.c
index bab9975755b..f144f39b083 100644
--- a/src/gallium/auxiliary/gallivm/lp_bld_nir.c
+++ b/src/gallium/auxiliary/gallivm/lp_bld_nir.c
@@ -1690,9 +1690,9 @@ static void visit_intrinsic(struct lp_build_nir_context *bld_base,
    case nir_intrinsic_load_base_instance:
    case nir_intrinsic_load_base_vertex:
    case nir_intrinsic_load_first_vertex:
-   case nir_intrinsic_load_work_group_id:
+   case nir_intrinsic_load_workgroup_id:
    case nir_intrinsic_load_local_invocation_id:
-   case nir_intrinsic_load_num_work_groups:
+   case nir_intrinsic_load_num_workgroups:
    case nir_intrinsic_load_invocation_id:
    case nir_intrinsic_load_front_face:
    case nir_intrinsic_load_draw_id:
diff --git a/src/gallium/auxiliary/gallivm/lp_bld_nir_soa.c b/src/gallium/auxiliary/gallivm/lp_bld_nir_soa.c
index 45f99571b21..ddd706cc37c 100644
--- a/src/gallium/auxiliary/gallivm/lp_bld_nir_soa.c
+++ b/src/gallium/auxiliary/gallivm/lp_bld_nir_soa.c
@@ -1514,7 +1514,7 @@ static void emit_sysval_intrin(struct lp_build_nir_context *bld_base,
    case nir_intrinsic_load_primitive_id:
       result[0] = bld->system_values.prim_id;
       break;
-   case nir_intrinsic_load_work_group_id: {
+   case nir_intrinsic_load_workgroup_id: {
       LLVMValueRef tmp[3];
       for (unsigned i = 0; i < 3; i++) {
          tmp[i] = LLVMBuildExtractElement(gallivm->builder, bld->system_values.block_id, lp_build_const_int32(gallivm, i), "");
@@ -1528,7 +1528,7 @@ static void emit_sysval_intrin(struct lp_build_nir_context *bld_base,
       for (unsigned i = 0; i < 3; i++)
          result[i] = LLVMBuildExtractValue(gallivm->builder, bld->system_values.thread_id, i, "");
       break;
-   case nir_intrinsic_load_num_work_groups: {
+   case nir_intrinsic_load_num_workgroups: {
       LLVMValueRef tmp[3];
       for (unsigned i = 0; i < 3; i++) {
          tmp[i] = LLVMBuildExtractElement(gallivm->builder, bld->system_values.grid_size, lp_build_const_int32(gallivm, i), "");
diff --git a/src/gallium/auxiliary/nir/nir_to_tgsi.c b/src/gallium/auxiliary/nir/nir_to_tgsi.c
index 69fd597887f..409e9188be8 100644
--- a/src/gallium/auxiliary/nir/nir_to_tgsi.c
+++ b/src/gallium/auxiliary/nir/nir_to_tgsi.c
@@ -1650,8 +1650,8 @@ ntt_emit_intrinsic(struct ntt_compile *c, nir_intrinsic_instr *instr)
    case nir_intrinsic_load_tess_level_outer:
    case nir_intrinsic_load_tess_level_inner:
    case nir_intrinsic_load_local_invocation_id:
-   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_subgroup_size:
    case nir_intrinsic_load_subgroup_invocation:
diff --git a/src/gallium/auxiliary/nir/nir_to_tgsi_info.c b/src/gallium/auxiliary/nir/nir_to_tgsi_info.c
index 8b1a9e76c18..aa9ce92dbae 100644
--- a/src/gallium/auxiliary/nir/nir_to_tgsi_info.c
+++ b/src/gallium/auxiliary/nir/nir_to_tgsi_info.c
@@ -220,7 +220,7 @@ static void scan_instruction(const struct nir_shader *nir,
       case nir_intrinsic_load_invocation_id:
          info->uses_invocationid = true;
          break;
-      case nir_intrinsic_load_num_work_groups:
+      case nir_intrinsic_load_num_workgroups:
          info->uses_grid_size = true;
          break;
       case nir_intrinsic_load_workgroup_size:
@@ -229,12 +229,12 @@ static void scan_instruction(const struct nir_shader *nir,
             info->uses_block_size = true;
          break;
       case nir_intrinsic_load_local_invocation_id:
-      case nir_intrinsic_load_work_group_id: {
+      case nir_intrinsic_load_workgroup_id: {
          unsigned mask = nir_ssa_def_components_read(&intr->dest.ssa);
          while (mask) {
             unsigned i = u_bit_scan(&mask);
 
-            if (intr->intrinsic == nir_intrinsic_load_work_group_id)
+            if (intr->intrinsic == nir_intrinsic_load_workgroup_id)
                info->uses_block_id[i] = true;
             else
                info->uses_thread_id[i] = true;
diff --git a/src/gallium/auxiliary/nir/tgsi_to_nir.c b/src/gallium/auxiliary/nir/tgsi_to_nir.c
index 339d1ccf6a1..5b625ef66a1 100644
--- a/src/gallium/auxiliary/nir/tgsi_to_nir.c
+++ b/src/gallium/auxiliary/nir/tgsi_to_nir.c
@@ -620,7 +620,7 @@ ttn_src_for_file_and_index(struct ttn_compile *c, unsigned file, unsigned index,
          load = nir_load_local_invocation_id(b);
          break;
       case TGSI_SEMANTIC_BLOCK_ID:
-         load = nir_load_work_group_id(b, 32);
+         load = nir_load_workgroup_id(b, 32);
          break;
       case TGSI_SEMANTIC_BLOCK_SIZE:
          load = nir_load_workgroup_size(b);
diff --git a/src/gallium/auxiliary/tgsi/tgsi_from_mesa.c b/src/gallium/auxiliary/tgsi/tgsi_from_mesa.c
index 88ee6bb6543..9e39497114c 100644
--- a/src/gallium/auxiliary/tgsi/tgsi_from_mesa.c
+++ b/src/gallium/auxiliary/tgsi/tgsi_from_mesa.c
@@ -265,9 +265,9 @@ tgsi_get_sysval_semantic(unsigned sysval)
    /* Compute shader */
    case SYSTEM_VALUE_LOCAL_INVOCATION_ID:
       return TGSI_SEMANTIC_THREAD_ID;
-   case SYSTEM_VALUE_WORK_GROUP_ID:
+   case SYSTEM_VALUE_WORKGROUP_ID:
       return TGSI_SEMANTIC_BLOCK_ID;
-   case SYSTEM_VALUE_NUM_WORK_GROUPS:
+   case SYSTEM_VALUE_NUM_WORKGROUPS:
       return TGSI_SEMANTIC_GRID_SIZE;
    case SYSTEM_VALUE_WORKGROUP_SIZE:
       return TGSI_SEMANTIC_BLOCK_SIZE;
diff --git a/src/gallium/drivers/freedreno/a5xx/fd5_compute.c b/src/gallium/drivers/freedreno/a5xx/fd5_compute.c
index 4f608473930..734221433e8 100644
--- a/src/gallium/drivers/freedreno/a5xx/fd5_compute.c
+++ b/src/gallium/drivers/freedreno/a5xx/fd5_compute.c
@@ -91,7 +91,7 @@ cs_program_emit(struct fd_ringbuffer *ring, struct ir3_shader_variant *v)
    uint32_t local_invocation_id, work_group_id;
    local_invocation_id =
       ir3_find_sysval_regid(v, SYSTEM_VALUE_LOCAL_INVOCATION_ID);
-   work_group_id = ir3_find_sysval_regid(v, SYSTEM_VALUE_WORK_GROUP_ID);
+   work_group_id = ir3_find_sysval_regid(v, SYSTEM_VALUE_WORKGROUP_ID);
 
    OUT_PKT4(ring, REG_A5XX_HLSQ_CS_CNTL_0, 2);
    OUT_RING(ring, A5XX_HLSQ_CS_CNTL_0_WGIDCONSTID(work_group_id) |
diff --git a/src/gallium/drivers/freedreno/a6xx/fd6_compute.c b/src/gallium/drivers/freedreno/a6xx/fd6_compute.c
index cbb3bed8e7f..bcb7433c754 100644
--- a/src/gallium/drivers/freedreno/a6xx/fd6_compute.c
+++ b/src/gallium/drivers/freedreno/a6xx/fd6_compute.c
@@ -78,7 +78,7 @@ cs_program_emit(struct fd_context *ctx, struct fd_ringbuffer *ring,
    uint32_t local_invocation_id, work_group_id;
    local_invocation_id =
       ir3_find_sysval_regid(v, SYSTEM_VALUE_LOCAL_INVOCATION_ID);
-   work_group_id = ir3_find_sysval_regid(v, SYSTEM_VALUE_WORK_GROUP_ID);
+   work_group_id = ir3_find_sysval_regid(v, SYSTEM_VALUE_WORKGROUP_ID);
 
    OUT_PKT4(ring, REG_A6XX_HLSQ_CS_CNTL_0, 2);
    OUT_RING(ring, A6XX_HLSQ_CS_CNTL_0_WGIDCONSTID(work_group_id) |
diff --git a/src/gallium/drivers/iris/iris_program.c b/src/gallium/drivers/iris/iris_program.c
index 9028085eb4b..2170e67267d 100644
--- a/src/gallium/drivers/iris/iris_program.c
+++ b/src/gallium/drivers/iris/iris_program.c
@@ -852,7 +852,7 @@ iris_setup_binding_table(const struct intel_device_info *devinfo,
 
          nir_intrinsic_instr *intrin = nir_instr_as_intrinsic(instr);
          switch (intrin->intrinsic) {
-         case nir_intrinsic_load_num_work_groups:
+         case nir_intrinsic_load_num_workgroups:
             bt->used_mask[IRIS_SURFACE_GROUP_CS_WORK_GROUPS] = 1;
             break;
 
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 62c544b8c21..90150065f86 100644
--- a/src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp
+++ b/src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp
@@ -1570,7 +1570,7 @@ Converter::convert(nir_intrinsic_op intr)
       return SV_NTID;
    case nir_intrinsic_load_local_invocation_id:
       return SV_TID;
-   case nir_intrinsic_load_num_work_groups:
+   case nir_intrinsic_load_num_workgroups:
       return SV_NCTAID;
    case nir_intrinsic_load_patch_vertices_in:
       return SV_VERTEX_COUNT;
@@ -1602,7 +1602,7 @@ Converter::convert(nir_intrinsic_op intr)
       return SV_TESS_OUTER;
    case nir_intrinsic_load_vertex_id:
       return SV_VERTEX_ID;
-   case nir_intrinsic_load_work_group_id:
+   case nir_intrinsic_load_workgroup_id:
       return SV_CTAID;
    case nir_intrinsic_load_work_dim:
       return SV_WORK_DIM;
@@ -1845,7 +1845,7 @@ Converter::visit(nir_intrinsic_instr *insn)
    case nir_intrinsic_load_invocation_id:
    case nir_intrinsic_load_workgroup_size:
    case nir_intrinsic_load_local_invocation_id:
-   case nir_intrinsic_load_num_work_groups:
+   case nir_intrinsic_load_num_workgroups:
    case nir_intrinsic_load_patch_vertices_in:
    case nir_intrinsic_load_primitive_id:
    case nir_intrinsic_load_sample_id:
@@ -1861,7 +1861,7 @@ Converter::visit(nir_intrinsic_instr *insn)
    case nir_intrinsic_load_tess_level_inner:
    case nir_intrinsic_load_tess_level_outer:
    case nir_intrinsic_load_vertex_id:
-   case nir_intrinsic_load_work_group_id:
+   case nir_intrinsic_load_workgroup_id:
    case nir_intrinsic_load_work_dim: {
       const DataType dType = getDType(insn);
       SVSemantic sv = convert(op);
diff --git a/src/gallium/drivers/r600/sfn/sfn_shader_compute.cpp b/src/gallium/drivers/r600/sfn/sfn_shader_compute.cpp
index 26ac54981af..7abb4144095 100644
--- a/src/gallium/drivers/r600/sfn/sfn_shader_compute.cpp
+++ b/src/gallium/drivers/r600/sfn/sfn_shader_compute.cpp
@@ -69,10 +69,10 @@ bool ComputeShaderFromNir::emit_intrinsic_instruction_override(nir_intrinsic_ins
    switch (instr->intrinsic) {
    case nir_intrinsic_load_local_invocation_id:
       return emit_load_3vec(instr, m_local_invocation_id);
-   case nir_intrinsic_load_work_group_id:
+   case nir_intrinsic_load_workgroup_id:
       return emit_load_3vec(instr, m_workgroup_id);
-   case nir_intrinsic_load_num_work_groups:
-      return emit_load_num_work_groups(instr);
+   case nir_intrinsic_load_num_workgroups:
+      return emit_load_num_workgroups(instr);
    default:
       return false;
    }
@@ -86,7 +86,7 @@ bool ComputeShaderFromNir::emit_load_3vec(nir_intrinsic_instr* instr,
    return true;
 }
 
-bool ComputeShaderFromNir::emit_load_num_work_groups(nir_intrinsic_instr* instr)
+bool ComputeShaderFromNir::emit_load_num_workgroups(nir_intrinsic_instr* instr)
 {
    PValue a_zero = get_temp_register(1);
    emit_instruction(new AluInstruction(op1_mov, a_zero, Value::zero, EmitInstruction::last_write));
diff --git a/src/gallium/drivers/r600/sfn/sfn_shader_compute.h b/src/gallium/drivers/r600/sfn/sfn_shader_compute.h
index fea6f0122a1..b70cbad754a 100644
--- a/src/gallium/drivers/r600/sfn/sfn_shader_compute.h
+++ b/src/gallium/drivers/r600/sfn/sfn_shader_compute.h
@@ -50,7 +50,7 @@ private:
    void do_finalize() override;
 
    bool emit_load_3vec(nir_intrinsic_instr* instr, const std::array<PValue,3>& src);
-   bool emit_load_num_work_groups(nir_intrinsic_instr* instr);
+   bool emit_load_num_workgroups(nir_intrinsic_instr* instr);
 
    int m_reserved_registers;
    std::array<PValue,3> m_workgroup_id;
diff --git a/src/gallium/drivers/radeonsi/si_shader_nir.c b/src/gallium/drivers/radeonsi/si_shader_nir.c
index aeeb0988e87..199f6528275 100644
--- a/src/gallium/drivers/radeonsi/si_shader_nir.c
+++ b/src/gallium/drivers/radeonsi/si_shader_nir.c
@@ -301,12 +301,12 @@ static void scan_instruction(const struct nir_shader *nir, struct si_shader_info
             info->uses_indirect_descriptor = true;
          break;
       case nir_intrinsic_load_local_invocation_id:
-      case nir_intrinsic_load_work_group_id: {
+      case nir_intrinsic_load_workgroup_id: {
          unsigned mask = nir_ssa_def_components_read(&intr->dest.ssa);
          while (mask) {
             unsigned i = u_bit_scan(&mask);
 
-            if (intr->intrinsic == nir_intrinsic_load_work_group_id)
+            if (intr->intrinsic == nir_intrinsic_load_workgroup_id)
                info->uses_block_id[i] = true;
             else
                info->uses_thread_id[i] = true;
@@ -431,7 +431,7 @@ void si_nir_scan_shader(const struct nir_shader *nir, struct si_shader_info *inf
    info->uses_base_vertex = BITSET_TEST(nir->info.system_values_read, SYSTEM_VALUE_BASE_VERTEX);
    info->uses_base_instance = BITSET_TEST(nir->info.system_values_read, SYSTEM_VALUE_BASE_INSTANCE);
    info->uses_invocationid = BITSET_TEST(nir->info.system_values_read, SYSTEM_VALUE_INVOCATION_ID);
-   info->uses_grid_size = BITSET_TEST(nir->info.system_values_read, SYSTEM_VALUE_NUM_WORK_GROUPS);
+   info->uses_grid_size = BITSET_TEST(nir->info.system_values_read, SYSTEM_VALUE_NUM_WORKGROUPS);
    info->uses_subgroup_info = BITSET_TEST(nir->info.system_values_read, SYSTEM_VALUE_LOCAL_INVOCATION_INDEX) ||
                               BITSET_TEST(nir->info.system_values_read, SYSTEM_VALUE_SUBGROUP_ID) ||
                               BITSET_TEST(nir->info.system_values_read, SYSTEM_VALUE_NUM_SUBGROUPS);
diff --git a/src/gallium/drivers/radeonsi/si_shaderlib_nir.c b/src/gallium/drivers/radeonsi/si_shaderlib_nir.c
index 8b9fd355cfc..c55985468b4 100644
--- a/src/gallium/drivers/radeonsi/si_shaderlib_nir.c
+++ b/src/gallium/drivers/radeonsi/si_shaderlib_nir.c
@@ -42,7 +42,7 @@ static nir_ssa_def *get_global_ids(nir_builder *b, unsigned num_components)
    unsigned mask = BITFIELD_MASK(num_components);
 
    nir_ssa_def *local_ids = nir_channels(b, nir_load_local_invocation_id(b), mask);
-   nir_ssa_def *block_ids = nir_channels(b, nir_load_work_group_id(b, 32), mask);
+   nir_ssa_def *block_ids = nir_channels(b, nir_load_workgroup_id(b, 32), mask);
    nir_ssa_def *block_size = nir_channels(b, nir_load_workgroup_size(b), mask);
    return nir_iadd(b, nir_imul(b, block_ids, block_size), local_ids);
 }
diff --git a/src/gallium/drivers/zink/nir_to_spirv/nir_to_spirv.c b/src/gallium/drivers/zink/nir_to_spirv/nir_to_spirv.c
index ab7eaa252a4..aa115f434cb 100644
--- a/src/gallium/drivers/zink/nir_to_spirv/nir_to_spirv.c
+++ b/src/gallium/drivers/zink/nir_to_spirv/nir_to_spirv.c
@@ -2765,11 +2765,11 @@ emit_intrinsic(struct ntv_context *ctx, nir_intrinsic_instr *intr)
       emit_image_intrinsic(ctx, intr);
       break;
 
-   case nir_intrinsic_load_work_group_id:
+   case nir_intrinsic_load_workgroup_id:
       emit_load_vec_input(ctx, intr, &ctx->workgroup_id_var, "gl_WorkGroupID", SpvBuiltInWorkgroupId, nir_type_uint);
       break;
 
-   case nir_intrinsic_load_num_work_groups:
+   case nir_intrinsic_load_num_workgroups:
       emit_load_vec_input(ctx, intr, &ctx->num_workgroups_var, "gl_NumWorkGroups", SpvBuiltInNumWorkgroups, nir_type_uint);
       break;
 
diff --git a/src/intel/compiler/brw_fs_nir.cpp b/src/intel/compiler/brw_fs_nir.cpp
index d50b7ab9f5f..4110cf0b621 100644
--- a/src/intel/compiler/brw_fs_nir.cpp
+++ b/src/intel/compiler/brw_fs_nir.cpp
@@ -191,10 +191,10 @@ emit_system_values_block(nir_block *block, fs_visitor *v)
             *reg = *v->emit_samplemaskin_setup();
          break;
 
-      case nir_intrinsic_load_work_group_id:
+      case nir_intrinsic_load_workgroup_id:
          assert(v->stage == MESA_SHADER_COMPUTE ||
                 v->stage == MESA_SHADER_KERNEL);
-         reg = &v->nir_system_values[SYSTEM_VALUE_WORK_GROUP_ID];
+         reg = &v->nir_system_values[SYSTEM_VALUE_WORKGROUP_ID];
          if (reg->file == BAD_FILE)
             *reg = *v->emit_cs_work_group_id_setup();
          break;
@@ -3691,7 +3691,7 @@ fs_visitor::nir_emit_cs_intrinsic(const fs_builder &bld,
       break;
 
    case nir_intrinsic_load_local_invocation_id:
-   case nir_intrinsic_load_work_group_id: {
+   case nir_intrinsic_load_workgroup_id: {
       gl_system_value sv = nir_system_value_from_intrinsic(instr->intrinsic);
       fs_reg val = nir_system_values[sv];
       assert(val.file != BAD_FILE);
@@ -3701,7 +3701,7 @@ fs_visitor::nir_emit_cs_intrinsic(const fs_builder &bld,
       break;
    }
 
-   case nir_intrinsic_load_num_work_groups: {
+   case nir_intrinsic_load_num_workgroups: {
       assert(nir_dest_bit_size(instr->dest) == 32);
       const unsigned surface =
          cs_prog_data->binding_table.work_groups_start;
diff --git a/src/intel/compiler/brw_nir_lower_cs_intrinsics.c b/src/intel/compiler/brw_nir_lower_cs_intrinsics.c
index 66999f41363..da5913e42e8 100644
--- a/src/intel/compiler/brw_nir_lower_cs_intrinsics.c
+++ b/src/intel/compiler/brw_nir_lower_cs_intrinsics.c
@@ -54,8 +54,8 @@ lower_cs_intrinsics_convert_block(struct lower_intrinsics_state *state,
       nir_ssa_def *sysval;
       switch (intrinsic->intrinsic) {
       case nir_intrinsic_load_workgroup_size:
-      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:
          /* Convert this to 32-bit if it's not */
          if (intrinsic->dest.ssa.bit_size == 64) {
             intrinsic->dest.ssa.bit_size = 32;
diff --git a/src/intel/compiler/brw_nir_rt.c b/src/intel/compiler/brw_nir_rt.c
index 9baf53e703d..93f0a8fc541 100644
--- a/src/intel/compiler/brw_nir_rt.c
+++ b/src/intel/compiler/brw_nir_rt.c
@@ -438,7 +438,7 @@ brw_nir_create_raygen_trampoline(const struct brw_compiler *compiler,
    nir_ssa_def *local_shift =
       nir_u2u32(&b, load_trampoline_param(&b, local_group_size_log2, 3, 8));
 
-   nir_ssa_def *global_id = nir_load_work_group_id(&b, 32);
+   nir_ssa_def *global_id = nir_load_workgroup_id(&b, 32);
    nir_ssa_def *simd_channel = nir_load_subgroup_invocation(&b);
    nir_ssa_def *local_x =
       nir_ubfe(&b, simd_channel, nir_imm_int(&b, 0),
diff --git a/src/intel/vulkan/anv_nir_add_base_work_group_id.c b/src/intel/vulkan/anv_nir_add_base_work_group_id.c
index c0435363688..97596214de9 100644
--- a/src/intel/vulkan/anv_nir_add_base_work_group_id.c
+++ b/src/intel/vulkan/anv_nir_add_base_work_group_id.c
@@ -44,7 +44,7 @@ anv_nir_add_base_work_group_id(nir_shader *shader)
                continue;
 
             nir_intrinsic_instr *load_id = nir_instr_as_intrinsic(instr);
-            if (load_id->intrinsic != nir_intrinsic_load_work_group_id)
+            if (load_id->intrinsic != nir_intrinsic_load_workgroup_id)
                continue;
 
             b.cursor = nir_after_instr(&load_id->instr);
diff --git a/src/microsoft/clc/clc_compiler.c b/src/microsoft/clc/clc_compiler.c
index 31e5f8c53dc..6eff749f3cc 100644
--- a/src/microsoft/clc/clc_compiler.c
+++ b/src/microsoft/clc/clc_compiler.c
@@ -351,10 +351,10 @@ clc_lower_64bit_semantics(nir_shader *nir)
                case nir_intrinsic_load_global_invocation_id_zero_base:
                case nir_intrinsic_load_base_global_invocation_id:
                case nir_intrinsic_load_local_invocation_id:
-               case nir_intrinsic_load_work_group_id:
-               case nir_intrinsic_load_work_group_id_zero_base:
-               case nir_intrinsic_load_base_work_group_id:
-               case nir_intrinsic_load_num_work_groups:
+               case nir_intrinsic_load_workgroup_id:
+               case nir_intrinsic_load_workgroup_id_zero_base:
+               case nir_intrinsic_load_base_workgroup_id:
+               case nir_intrinsic_load_num_workgroups:
                   break;
                default:
                   continue;
@@ -1322,7 +1322,7 @@ clc_to_dxil(struct clc_context *ctx,
 
    nir_lower_compute_system_values_options compute_options = {
       .has_base_global_invocation_id = (conf && conf->support_global_work_id_offsets),
-      .has_base_work_group_id = (conf && conf->support_work_group_id_offsets),
+      .has_base_workgroup_id = (conf && conf->support_workgroup_id_offsets),
    };
    NIR_PASS_V(nir, nir_lower_compute_system_values, &compute_options);
 
diff --git a/src/microsoft/clc/clc_compiler.h b/src/microsoft/clc/clc_compiler.h
index 6981bbaf415..fc772605b49 100644
--- a/src/microsoft/clc/clc_compiler.h
+++ b/src/microsoft/clc/clc_compiler.h
@@ -233,7 +233,7 @@ struct clc_runtime_kernel_conf {
    struct clc_runtime_arg_info *args;
    unsigned lower_bit_size;
    unsigned support_global_work_id_offsets;
-   unsigned support_work_group_id_offsets;
+   unsigned support_workgroup_id_offsets;
 };
 
 struct clc_dxil_object *
diff --git a/src/microsoft/clc/clc_nir.c b/src/microsoft/clc/clc_nir.c
index 3d8fb7c9772..fddff035c71 100644
--- a/src/microsoft/clc/clc_nir.c
+++ b/src/microsoft/clc/clc_nir.c
@@ -84,8 +84,8 @@ lower_load_local_group_size(nir_builder *b, nir_intrinsic_instr *intr)
 }
 
 static bool
-lower_load_num_work_groups(nir_builder *b, nir_intrinsic_instr *intr,
-                           nir_variable *var)
+lower_load_num_workgroups(nir_builder *b, nir_intrinsic_instr *intr,
+                          nir_variable *var)
 {
    b->cursor = nir_after_instr(&intr->instr);
 
@@ -102,7 +102,7 @@ lower_load_num_work_groups(nir_builder *b, nir_intrinsic_instr *intr,
 }
 
 static bool
-lower_load_base_work_group_id(nir_builder *b, nir_intrinsic_instr *intr,
+lower_load_base_workgroup_id(nir_builder *b, nir_intrinsic_instr *intr,
                              nir_variable *var)
 {
    b->cursor = nir_after_instr(&intr->instr);
@@ -149,11 +149,11 @@ clc_nir_lower_system_values(nir_shader *nir, nir_variable *var)
             case nir_intrinsic_load_workgroup_size:
                lower_load_local_group_size(&b, intr);
                break;
-            case nir_intrinsic_load_num_work_groups:
-               lower_load_num_work_groups(&b, intr, var);
+            case nir_intrinsic_load_num_workgroups:
+               lower_load_num_workgroups(&b, intr, var);
                break;
-            case nir_intrinsic_load_base_work_group_id:
-               lower_load_base_work_group_id(&b, intr, var);
+            case nir_intrinsic_load_base_workgroup_id:
+               lower_load_base_workgroup_id(&b, intr, var);
                break;
             default: break;
             }
diff --git a/src/microsoft/clc/compute_test.cpp b/src/microsoft/clc/compute_test.cpp
index 2854268267c..d19ff21095b 100644
--- a/src/microsoft/clc/compute_test.cpp
+++ b/src/microsoft/clc/compute_test.cpp
@@ -467,7 +467,7 @@ ComputeTest::run_shader_with_raw_args(Shader shader,
       compile_args.work_props.global_offset_x != 0 ||
       compile_args.work_props.global_offset_y != 0 ||
       compile_args.work_props.global_offset_z != 0;
-   conf.support_work_group_id_offsets =
+   conf.support_workgroup_id_offsets =
       compile_args.work_props.group_id_offset_x != 0 ||
       compile_args.work_props.group_id_offset_y != 0 ||
       compile_args.work_props.group_id_offset_z != 0;
@@ -493,7 +493,7 @@ ComputeTest::run_shader_with_raw_args(Shader shader,
    std::vector<uint8_t> argsbuf(dxil->metadata.kernel_inputs_buf_size);
    std::vector<ComPtr<ID3D12Resource>> argres(shader.dxil->kernel->num_args);
    clc_work_properties_data work_props = compile_args.work_props;
-   if (!conf.support_work_group_id_offsets) {
+   if (!conf.support_workgroup_id_offsets) {
       work_props.group_count_total_x = compile_args.x / conf.local_size[0];
       work_props.group_count_total_y = compile_args.y / conf.local_size[1];
       work_props.group_count_total_z = compile_args.z / conf.local_size[2];
diff --git a/src/microsoft/compiler/nir_to_dxil.c b/src/microsoft/compiler/nir_to_dxil.c
index 6c16268f09d..9952c395e52 100644
--- a/src/microsoft/compiler/nir_to_dxil.c
+++ b/src/microsoft/compiler/nir_to_dxil.c
@@ -2264,7 +2264,7 @@ emit_load_local_invocation_id(struct ntd_context *ctx,
 }
 
 static bool
-emit_load_local_work_group_id(struct ntd_context *ctx,
+emit_load_local_workgroup_id(struct ntd_context *ctx,
                               nir_intrinsic_instr *intr)
 {
    assert(intr->dest.is_ssa);
@@ -3354,9 +3354,9 @@ emit_intrinsic(struct ntd_context *ctx, nir_intrinsic_instr *intr)
       return emit_load_global_invocation_id(ctx, intr);
    case nir_intrinsic_load_local_invocation_id:
       return emit_load_local_invocation_id(ctx, intr);
-   case nir_intrinsic_load_work_group_id:
-   case nir_intrinsic_load_work_group_id_zero_base:
-      return emit_load_local_work_group_id(ctx, intr);
+   case nir_intrinsic_load_workgroup_id:
+   case nir_intrinsic_load_workgroup_id_zero_base:
+      return emit_load_local_workgroup_id(ctx, intr);
    case nir_intrinsic_load_ssbo:
       return emit_load_ssbo(ctx, intr);
    case nir_intrinsic_store_ssbo:
@@ -3455,7 +3455,7 @@ emit_intrinsic(struct ntd_context *ctx, nir_intrinsic_instr *intr)
    case nir_intrinsic_load_vulkan_descriptor:
       return emit_load_vulkan_descriptor(ctx, intr);
 
-   case nir_intrinsic_load_num_work_groups:
+   case nir_intrinsic_load_num_workgroups:
    case nir_intrinsic_load_workgroup_size:
    default:
       NIR_INSTR_UNSUPPORTED(&intr->instr);
diff --git a/src/panfrost/bifrost/bifrost_compile.c b/src/panfrost/bifrost/bifrost_compile.c
index 9189618e187..1f6199f639f 100644
--- a/src/panfrost/bifrost/bifrost_compile.c
+++ b/src/panfrost/bifrost/bifrost_compile.c
@@ -1176,7 +1176,7 @@ bi_emit_intrinsic(bi_builder *b, nir_intrinsic_instr *instr)
 
         case nir_intrinsic_load_viewport_scale:
         case nir_intrinsic_load_viewport_offset:
-        case nir_intrinsic_load_num_work_groups:
+        case nir_intrinsic_load_num_workgroups:
         case nir_intrinsic_load_workgroup_size:
                 bi_load_sysval_nir(b, instr, 3, 0);
                 break;
@@ -1238,7 +1238,7 @@ bi_emit_intrinsic(bi_builder *b, nir_intrinsic_instr *instr)
                                          bi_half(bi_register(55 + i / 2), i % 2));
                 break;
 
-        case nir_intrinsic_load_work_group_id:
+        case nir_intrinsic_load_workgroup_id:
                 for (unsigned i = 0; i < 3; ++i)
                         bi_mov_i32_to(b, bi_word(dst, i), bi_register(57 + i));
                 break;
diff --git a/src/panfrost/midgard/midgard_compile.c b/src/panfrost/midgard/midgard_compile.c
index d0ae5a79c19..74b185b506e 100644
--- a/src/panfrost/midgard/midgard_compile.c
+++ b/src/panfrost/midgard/midgard_compile.c
@@ -1515,7 +1515,7 @@ static unsigned
 compute_builtin_arg(nir_intrinsic_op op)
 {
         switch (op) {
-        case nir_intrinsic_load_work_group_id:
+        case nir_intrinsic_load_workgroup_id:
                 return REGISTER_LDST_GROUP_ID;
         case nir_intrinsic_load_local_invocation_id:
                 return REGISTER_LDST_LOCAL_THREAD_ID;
@@ -2023,13 +2023,13 @@ emit_intrinsic(compiler_context *ctx, nir_intrinsic_instr *instr)
 
         case nir_intrinsic_load_viewport_scale:
         case nir_intrinsic_load_viewport_offset:
-        case nir_intrinsic_load_num_work_groups:
+        case nir_intrinsic_load_num_workgroups:
         case nir_intrinsic_load_sampler_lod_parameters_pan:
         case nir_intrinsic_load_workgroup_size:
                 emit_sysval_read(ctx, &instr->instr, 3, 0);
                 break;
 
-        case nir_intrinsic_load_work_group_id:
+        case nir_intrinsic_load_workgroup_id:
         case nir_intrinsic_load_local_invocation_id:
         case nir_intrinsic_load_global_invocation_id:
         case nir_intrinsic_load_global_invocation_id_zero_base:
diff --git a/src/panfrost/util/pan_lower_64bit_intrin.c b/src/panfrost/util/pan_lower_64bit_intrin.c
index 6dec682cc15..7c4edcfa9d7 100644
--- a/src/panfrost/util/pan_lower_64bit_intrin.c
+++ b/src/panfrost/util/pan_lower_64bit_intrin.c
@@ -44,8 +44,8 @@ nir_lower_64bit_intrin_instr(nir_builder *b, nir_instr *instr, void *data)
         switch (intr->intrinsic) {
         case nir_intrinsic_load_global_invocation_id:
         case nir_intrinsic_load_global_invocation_id_zero_base:
-        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:
                 break;
 
         default:
diff --git a/src/panfrost/util/pan_sysval.c b/src/panfrost/util/pan_sysval.c
index 0715058d326..80a509f5b8b 100644
--- a/src/panfrost/util/pan_sysval.c
+++ b/src/panfrost/util/pan_sysval.c
@@ -70,7 +70,7 @@ panfrost_nir_sysval_for_intrinsic(nir_intrinsic_instr *instr)
                 return PAN_SYSVAL_VIEWPORT_SCALE;
         case nir_intrinsic_load_viewport_offset:
                 return PAN_SYSVAL_VIEWPORT_OFFSET;
-        case nir_intrinsic_load_num_work_groups:
+        case nir_intrinsic_load_num_workgroups:
                 return PAN_SYSVAL_NUM_WORK_GROUPS;
         case nir_intrinsic_load_workgroup_size:
                 return PAN_SYSVAL_LOCAL_GROUP_SIZE;