nir: Rename nir_intrinsic_load_local_group_size to nir_intrinsic_load_workgroup_size

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>

diff --git a/src/amd/llvm/ac_nir_to_llvm.c b/src/amd/llvm/ac_nir_to_llvm.c
index be41ced..5fabf94 100644 (file)
--- a/src/amd/llvm/ac_nir_to_llvm.c
+++ b/src/amd/llvm/ac_nir_to_llvm.c
@@ -3417,7 +3417,7 @@ static void visit_intrinsic(struct ac_nir_context *ctx, nir_intrinsic_instr *ins
       result = ctx->abi->load_base_vertex(ctx->abi,
                                           instr->intrinsic == nir_intrinsic_load_base_vertex);
       break;
-   case nir_intrinsic_load_local_group_size:
+   case nir_intrinsic_load_workgroup_size:
       result = ctx->abi->load_local_group_size(ctx->abi);
       break;
    case nir_intrinsic_load_vertex_id:

diff --git a/src/compiler/nir/nir.c b/src/compiler/nir/nir.c
index 9aa91cd..2a78308 100644 (file)
--- a/src/compiler/nir/nir.c
+++ b/src/compiler/nir/nir.c
@@ -2020,7 +2020,7 @@ nir_intrinsic_from_system_value(gl_system_value val)
    case SYSTEM_VALUE_SUBGROUP_ID:
       return nir_intrinsic_load_subgroup_id;
    case SYSTEM_VALUE_WORKGROUP_SIZE:
-      return nir_intrinsic_load_local_group_size;
+      return nir_intrinsic_load_workgroup_size;
    case SYSTEM_VALUE_GLOBAL_INVOCATION_ID:
       return nir_intrinsic_load_global_invocation_id;
    case SYSTEM_VALUE_BASE_GLOBAL_INVOCATION_ID:
@@ -2150,7 +2150,7 @@ nir_system_value_from_intrinsic(nir_intrinsic_op intrin)
       return SYSTEM_VALUE_NUM_SUBGROUPS;
    case nir_intrinsic_load_subgroup_id:
       return SYSTEM_VALUE_SUBGROUP_ID;
-   case nir_intrinsic_load_local_group_size:
+   case nir_intrinsic_load_workgroup_size:
       return SYSTEM_VALUE_WORKGROUP_SIZE;
    case nir_intrinsic_load_global_invocation_id:
       return SYSTEM_VALUE_GLOBAL_INVOCATION_ID;

diff --git a/src/compiler/nir/nir_divergence_analysis.c b/src/compiler/nir/nir_divergence_analysis.c
index 9777ee7..05d525f 100644 (file)
--- a/src/compiler/nir/nir_divergence_analysis.c
+++ b/src/compiler/nir/nir_divergence_analysis.c
@@ -102,7 +102,7 @@ visit_intrinsic(nir_shader *shader, nir_intrinsic_instr *instr)
    case nir_intrinsic_load_push_constant:
    case nir_intrinsic_load_work_dim:
    case nir_intrinsic_load_num_work_groups:
-   case nir_intrinsic_load_local_group_size:
+   case nir_intrinsic_load_workgroup_size:
    case nir_intrinsic_load_subgroup_id:
    case nir_intrinsic_load_num_subgroups:
    case nir_intrinsic_load_subgroup_size:

diff --git a/src/compiler/nir/nir_gather_info.c b/src/compiler/nir/nir_gather_info.c
index 5e8ad30..bdf84a3 100644 (file)
--- a/src/compiler/nir/nir_gather_info.c
+++ b/src/compiler/nir/nir_gather_info.c
@@ -625,7 +625,7 @@ gather_intrinsic_info(nir_intrinsic_instr *instr, nir_shader *shader,
    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_local_group_size:
+   case nir_intrinsic_load_workgroup_size:
    case nir_intrinsic_load_work_dim:
    case nir_intrinsic_load_user_data_amd:
    case nir_intrinsic_load_view_index:

diff --git a/src/compiler/nir/nir_intrinsics.py b/src/compiler/nir/nir_intrinsics.py
index 18d24b9..89f8992 100644 (file)
--- a/src/compiler/nir/nir_intrinsics.py
+++ b/src/compiler/nir/nir_intrinsics.py
@@ -702,9 +702,9 @@ system_value("subgroup_le_mask", 0, bit_sizes=[32, 64])
 system_value("subgroup_lt_mask", 0, bit_sizes=[32, 64])
 system_value("num_subgroups", 1)
 system_value("subgroup_id", 1)
-system_value("local_group_size", 3)
+system_value("workgroup_size", 3)
 # note: the definition of global_invocation_id_zero_base is based on
-# (work_group_id * local_group_size) + local_invocation_id.
+# (work_group_id * workgroup_size) + local_invocation_id.
 # it is *not* based on work_group_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])

diff --git a/src/compiler/nir/nir_lower_system_values.c b/src/compiler/nir/nir_lower_system_values.c
index 20e9603..5db6c2b 100644 (file)
--- a/src/compiler/nir/nir_lower_system_values.c
+++ b/src/compiler/nir/nir_lower_system_values.c
@@ -54,7 +54,7 @@ sanitize_32bit_sysval(nir_builder *b, nir_intrinsic_instr *intrin)
 static nir_ssa_def*
 build_global_group_size(nir_builder *b, unsigned bit_size)
 {
-   nir_ssa_def *group_size = nir_load_local_group_size(b);
+   nir_ssa_def *group_size = nir_load_workgroup_size(b);
    nir_ssa_def *num_work_groups = nir_load_num_work_groups(b, bit_size);
    return nir_imul(b, nir_u2u(b, group_size, bit_size),
                       num_work_groups);
@@ -116,7 +116,7 @@ lower_system_value_instr(nir_builder *b, nir_instr *instr, void *_state)
 
    case nir_intrinsic_load_local_invocation_id:
    case nir_intrinsic_load_local_invocation_index:
-   case nir_intrinsic_load_local_group_size:
+   case nir_intrinsic_load_workgroup_size:
       return sanitize_32bit_sysval(b, intrin);
 
    case nir_intrinsic_load_deref: {
@@ -294,7 +294,7 @@ lower_compute_system_value_instr(nir_builder *b,
           * large so it can safely be omitted.
           */
          nir_ssa_def *local_index = nir_load_local_invocation_index(b);
-         nir_ssa_def *local_size = nir_load_local_group_size(b);
+         nir_ssa_def *local_size = nir_load_workgroup_size(b);
 
          /* Because no hardware supports a local workgroup size greater than
           * about 1K, this calculation can be done in 32-bit and can save some
@@ -324,7 +324,7 @@ lower_compute_system_value_instr(nir_builder *b,
          nir_ssa_def *size_x_imm;
 
          if (b->shader->info.cs.workgroup_size_variable)
-            size_x_imm = nir_channel(b, nir_load_local_group_size(b), 0);
+            size_x_imm = nir_channel(b, nir_load_workgroup_size(b), 0);
          else
             size_x_imm = nir_imm_int(b, size_x);
 
@@ -424,7 +424,7 @@ lower_compute_system_value_instr(nir_builder *b,
          return NULL;
       }
 
-   case nir_intrinsic_load_local_group_size:
+   case nir_intrinsic_load_workgroup_size:
       if (b->shader->info.cs.workgroup_size_variable) {
          /* If the local work group size is variable it can't be lowered at
           * this point.  We do, however, have to make sure that the intrinsic
@@ -445,7 +445,7 @@ 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) ||
           !b->shader->options->has_cs_global_id) {
-         nir_ssa_def *group_size = nir_load_local_group_size(b);
+         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 *local_id = nir_load_local_invocation_id(b);
 

diff --git a/src/freedreno/ir3/ir3_compiler_nir.c b/src/freedreno/ir3/ir3_compiler_nir.c
index 195b38f..1ff3aa3 100644 (file)
--- a/src/freedreno/ir3/ir3_compiler_nir.c
+++ b/src/freedreno/ir3/ir3_compiler_nir.c
@@ -2042,7 +2042,7 @@ emit_intrinsic(struct ir3_context *ctx, nir_intrinsic_instr *intr)
                        dst[i] = create_driver_param(ctx, IR3_DP_NUM_WORK_GROUPS_X + i);
                }
                break;
-       case nir_intrinsic_load_local_group_size:
+       case nir_intrinsic_load_workgroup_size:
                for (int i = 0; i < dest_components; i++) {
                        dst[i] = create_driver_param(ctx, IR3_DP_LOCAL_GROUP_SIZE_X + i);
                }

diff --git a/src/freedreno/ir3/ir3_nir.c b/src/freedreno/ir3/ir3_nir.c
index c8d0a0b..0353bb3 100644 (file)
--- a/src/freedreno/ir3/ir3_nir.c
+++ b/src/freedreno/ir3/ir3_nir.c
@@ -680,7 +680,7 @@ ir3_nir_scan_driver_consts(nir_shader *shader,
                                        layout->num_driver_params =
                                                MAX2(layout->num_driver_params, IR3_DP_NUM_WORK_GROUPS_Z + 1);
                                        break;
-                               case nir_intrinsic_load_local_group_size:
+                               case nir_intrinsic_load_workgroup_size:
                                        layout->num_driver_params =
                                                MAX2(layout->num_driver_params, IR3_DP_LOCAL_GROUP_SIZE_Z + 1);
                                        break;

diff --git a/src/gallium/auxiliary/gallivm/lp_bld_nir.c b/src/gallium/auxiliary/gallivm/lp_bld_nir.c
index 5d2c27e..bab9975 100644 (file)
--- a/src/gallium/auxiliary/gallivm/lp_bld_nir.c
+++ b/src/gallium/auxiliary/gallivm/lp_bld_nir.c
@@ -1696,7 +1696,7 @@ static void visit_intrinsic(struct lp_build_nir_context *bld_base,
    case nir_intrinsic_load_invocation_id:
    case nir_intrinsic_load_front_face:
    case nir_intrinsic_load_draw_id:
-   case nir_intrinsic_load_local_group_size:
+   case nir_intrinsic_load_workgroup_size:
    case nir_intrinsic_load_work_dim:
    case nir_intrinsic_load_tess_coord:
    case nir_intrinsic_load_tess_level_outer:

diff --git a/src/gallium/auxiliary/gallivm/lp_bld_nir_soa.c b/src/gallium/auxiliary/gallivm/lp_bld_nir_soa.c
index 05e5208..45f9957 100644 (file)
--- a/src/gallium/auxiliary/gallivm/lp_bld_nir_soa.c
+++ b/src/gallium/auxiliary/gallivm/lp_bld_nir_soa.c
@@ -1552,7 +1552,7 @@ static void emit_sysval_intrin(struct lp_build_nir_context *bld_base,
       break;
    default:
       break;
-   case nir_intrinsic_load_local_group_size:
+   case nir_intrinsic_load_workgroup_size:
      for (unsigned i = 0; i < 3; i++)
        result[i] = lp_build_broadcast_scalar(&bld_base->uint_bld, LLVMBuildExtractElement(gallivm->builder, bld->system_values.block_size, lp_build_const_int32(gallivm, i), ""));
      break;

diff --git a/src/gallium/auxiliary/nir/nir_to_tgsi.c b/src/gallium/auxiliary/nir/nir_to_tgsi.c
index 3c73d34..69fd597 100644 (file)
--- a/src/gallium/auxiliary/nir/nir_to_tgsi.c
+++ b/src/gallium/auxiliary/nir/nir_to_tgsi.c
@@ -1652,7 +1652,7 @@ ntt_emit_intrinsic(struct ntt_compile *c, nir_intrinsic_instr *instr)
    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_local_group_size:
+   case nir_intrinsic_load_workgroup_size:
    case nir_intrinsic_load_subgroup_size:
    case nir_intrinsic_load_subgroup_invocation:
    case nir_intrinsic_load_subgroup_eq_mask:

diff --git a/src/gallium/auxiliary/nir/nir_to_tgsi_info.c b/src/gallium/auxiliary/nir/nir_to_tgsi_info.c
index e74e90b..8b1a9e7 100644 (file)
--- a/src/gallium/auxiliary/nir/nir_to_tgsi_info.c
+++ b/src/gallium/auxiliary/nir/nir_to_tgsi_info.c
@@ -223,7 +223,7 @@ static void scan_instruction(const struct nir_shader *nir,
       case nir_intrinsic_load_num_work_groups:
          info->uses_grid_size = true;
          break;
-      case nir_intrinsic_load_local_group_size:
+      case nir_intrinsic_load_workgroup_size:
          /* The block size is translated to IMM with a fixed block size. */
          if (info->properties[TGSI_PROPERTY_CS_FIXED_BLOCK_WIDTH] == 0)
             info->uses_block_size = true;

diff --git a/src/gallium/auxiliary/nir/tgsi_to_nir.c b/src/gallium/auxiliary/nir/tgsi_to_nir.c
index 0a7de87..339d1cc 100644 (file)
--- a/src/gallium/auxiliary/nir/tgsi_to_nir.c
+++ b/src/gallium/auxiliary/nir/tgsi_to_nir.c
@@ -623,7 +623,7 @@ ttn_src_for_file_and_index(struct ttn_compile *c, unsigned file, unsigned index,
          load = nir_load_work_group_id(b, 32);
          break;
       case TGSI_SEMANTIC_BLOCK_SIZE:
-         load = nir_load_local_group_size(b);
+         load = nir_load_workgroup_size(b);
          break;
       case TGSI_SEMANTIC_CS_USER_DATA_AMD:
          load = nir_load_user_data_amd(b);

diff --git a/src/gallium/drivers/iris/iris_program.c b/src/gallium/drivers/iris/iris_program.c
index 92915a9..9028085 100644 (file)
--- a/src/gallium/drivers/iris/iris_program.c
+++ b/src/gallium/drivers/iris/iris_program.c
@@ -527,7 +527,7 @@ iris_setup_uniforms(const struct brw_compiler *compiler,
                                nir_intrinsic_base(intrin) * 16));
             break;
          }
-         case nir_intrinsic_load_local_group_size: {
+         case nir_intrinsic_load_workgroup_size: {
             assert(nir->info.cs.workgroup_size_variable);
             if (variable_group_size_idx == -1) {
                variable_group_size_idx = num_system_values;

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 ad7d104..62c544b 100644 (file)
--- a/src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp
+++ b/src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp
@@ -1566,7 +1566,7 @@ Converter::convert(nir_intrinsic_op intr)
       return SV_INSTANCE_ID;
    case nir_intrinsic_load_invocation_id:
       return SV_INVOCATION_ID;
-   case nir_intrinsic_load_local_group_size:
+   case nir_intrinsic_load_workgroup_size:
       return SV_NTID;
    case nir_intrinsic_load_local_invocation_id:
       return SV_TID;
@@ -1843,7 +1843,7 @@ Converter::visit(nir_intrinsic_instr *insn)
    case nir_intrinsic_load_helper_invocation:
    case nir_intrinsic_load_instance_id:
    case nir_intrinsic_load_invocation_id:
-   case nir_intrinsic_load_local_group_size:
+   case nir_intrinsic_load_workgroup_size:
    case nir_intrinsic_load_local_invocation_id:
    case nir_intrinsic_load_num_work_groups:
    case nir_intrinsic_load_patch_vertices_in:

diff --git a/src/gallium/drivers/radeonsi/si_shaderlib_nir.c b/src/gallium/drivers/radeonsi/si_shaderlib_nir.c
index ab43b8e..8b9fd35 100644 (file)
--- a/src/gallium/drivers/radeonsi/si_shaderlib_nir.c
+++ b/src/gallium/drivers/radeonsi/si_shaderlib_nir.c
@@ -43,7 +43,7 @@ static nir_ssa_def *get_global_ids(nir_builder *b, unsigned 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_size = nir_channels(b, nir_load_local_group_size(b), 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 a248379..ab7eaa2 100644 (file)
--- a/src/gallium/drivers/zink/nir_to_spirv/nir_to_spirv.c
+++ b/src/gallium/drivers/zink/nir_to_spirv/nir_to_spirv.c
@@ -2785,7 +2785,7 @@ emit_intrinsic(struct ntv_context *ctx, nir_intrinsic_instr *intr)
       emit_load_uint_input(ctx, intr, &ctx->local_invocation_index_var, "gl_LocalInvocationIndex", SpvBuiltInLocalInvocationIndex);
       break;
 
-   case nir_intrinsic_load_local_group_size: {
+   case nir_intrinsic_load_workgroup_size: {
       assert(ctx->local_group_size_var);
       store_dest(ctx, &intr->dest, ctx->local_group_size_var, nir_type_uint);
       break;

diff --git a/src/intel/compiler/brw_fs_nir.cpp b/src/intel/compiler/brw_fs_nir.cpp
index eac4429..d50b7ab 100644 (file)
--- a/src/intel/compiler/brw_fs_nir.cpp
+++ b/src/intel/compiler/brw_fs_nir.cpp
@@ -3814,7 +3814,7 @@ fs_visitor::nir_emit_cs_intrinsic(const fs_builder &bld,
       break;
    }
 
-   case nir_intrinsic_load_local_group_size: {
+   case nir_intrinsic_load_workgroup_size: {
       assert(compiler->lower_variable_group_size);
       assert(nir->info.cs.workgroup_size_variable);
       for (unsigned i = 0; i < 3; i++) {

diff --git a/src/intel/compiler/brw_nir_lower_cs_intrinsics.c b/src/intel/compiler/brw_nir_lower_cs_intrinsics.c
index 5dc6dd7..66999f4 100644 (file)
--- a/src/intel/compiler/brw_nir_lower_cs_intrinsics.c
+++ b/src/intel/compiler/brw_nir_lower_cs_intrinsics.c
@@ -53,7 +53,7 @@ lower_cs_intrinsics_convert_block(struct lower_intrinsics_state *state,
 
       nir_ssa_def *sysval;
       switch (intrinsic->intrinsic) {
-      case nir_intrinsic_load_local_group_size:
+      case nir_intrinsic_load_workgroup_size:
       case nir_intrinsic_load_work_group_id:
       case nir_intrinsic_load_num_work_groups:
          /* Convert this to 32-bit if it's not */
@@ -82,7 +82,7 @@ lower_cs_intrinsics_convert_block(struct lower_intrinsics_state *state,
             nir_ssa_def *size_x;
             nir_ssa_def *size_y;
             if (state->nir->info.cs.workgroup_size_variable) {
-               nir_ssa_def *size_xyz = nir_load_local_group_size(b);
+               nir_ssa_def *size_xyz = nir_load_workgroup_size(b);
                size_x = nir_channel(b, size_xyz, 0);
                size_y = nir_channel(b, size_xyz, 1);
             } else {
@@ -214,7 +214,7 @@ lower_cs_intrinsics_convert_block(struct lower_intrinsics_state *state,
       case nir_intrinsic_load_num_subgroups: {
          nir_ssa_def *size;
          if (state->nir->info.cs.workgroup_size_variable) {
-            nir_ssa_def *size_xyz = nir_load_local_group_size(b);
+            nir_ssa_def *size_xyz = nir_load_workgroup_size(b);
             nir_ssa_def *size_x = nir_channel(b, size_xyz, 0);
             nir_ssa_def *size_y = nir_channel(b, size_xyz, 1);
             nir_ssa_def *size_z = nir_channel(b, size_xyz, 2);

diff --git a/src/microsoft/clc/clc_nir.c b/src/microsoft/clc/clc_nir.c
index eaac4e6..3d8fb7c 100644 (file)
--- a/src/microsoft/clc/clc_nir.c
+++ b/src/microsoft/clc/clc_nir.c
@@ -146,7 +146,7 @@ clc_nir_lower_system_values(nir_shader *nir, nir_variable *var)
             case nir_intrinsic_load_work_dim:
                progress |= lower_load_work_dim(&b, intr, var);
                break;
-            case nir_intrinsic_load_local_group_size:
+            case nir_intrinsic_load_workgroup_size:
                lower_load_local_group_size(&b, intr);
                break;
             case nir_intrinsic_load_num_work_groups:

diff --git a/src/microsoft/compiler/nir_to_dxil.c b/src/microsoft/compiler/nir_to_dxil.c
index 3ea4183..6c16268 100644 (file)
--- a/src/microsoft/compiler/nir_to_dxil.c
+++ b/src/microsoft/compiler/nir_to_dxil.c
@@ -3456,7 +3456,7 @@ emit_intrinsic(struct ntd_context *ctx, nir_intrinsic_instr *intr)
       return emit_load_vulkan_descriptor(ctx, intr);
 
    case nir_intrinsic_load_num_work_groups:
-   case nir_intrinsic_load_local_group_size:
+   case nir_intrinsic_load_workgroup_size:
    default:
       NIR_INSTR_UNSUPPORTED(&intr->instr);
       assert("Unimplemented intrinsic instruction");

diff --git a/src/panfrost/bifrost/bifrost_compile.c b/src/panfrost/bifrost/bifrost_compile.c
index d73987d..9189618 100644 (file)
--- a/src/panfrost/bifrost/bifrost_compile.c
+++ b/src/panfrost/bifrost/bifrost_compile.c
@@ -1177,7 +1177,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_local_group_size:
+        case nir_intrinsic_load_workgroup_size:
                 bi_load_sysval_nir(b, instr, 3, 0);
                 break;
 

diff --git a/src/panfrost/midgard/midgard_compile.c b/src/panfrost/midgard/midgard_compile.c
index c04d591..d0ae5a7 100644 (file)
--- a/src/panfrost/midgard/midgard_compile.c
+++ b/src/panfrost/midgard/midgard_compile.c
@@ -2025,7 +2025,7 @@ emit_intrinsic(compiler_context *ctx, nir_intrinsic_instr *instr)
         case nir_intrinsic_load_viewport_offset:
         case nir_intrinsic_load_num_work_groups:
         case nir_intrinsic_load_sampler_lod_parameters_pan:
-        case nir_intrinsic_load_local_group_size:
+        case nir_intrinsic_load_workgroup_size:
                 emit_sysval_read(ctx, &instr->instr, 3, 0);
                 break;
 

diff --git a/src/panfrost/util/pan_sysval.c b/src/panfrost/util/pan_sysval.c
index fd261a8..0715058 100644 (file)
--- a/src/panfrost/util/pan_sysval.c
+++ b/src/panfrost/util/pan_sysval.c
@@ -72,7 +72,7 @@ panfrost_nir_sysval_for_intrinsic(nir_intrinsic_instr *instr)
                 return PAN_SYSVAL_VIEWPORT_OFFSET;
         case nir_intrinsic_load_num_work_groups:
                 return PAN_SYSVAL_NUM_WORK_GROUPS;
-        case nir_intrinsic_load_local_group_size:
+        case nir_intrinsic_load_workgroup_size:
                 return PAN_SYSVAL_LOCAL_GROUP_SIZE;
         case nir_intrinsic_load_work_dim:
                 return PAN_SYSVAL_WORK_DIM;