diff --git a/src/intel/compiler/brw_kernel.c b/src/intel/compiler/brw_kernel.c
index ea0aa2d248c..8d00b2a706d 100644
--- a/src/intel/compiler/brw_kernel.c
+++ b/src/intel/compiler/brw_kernel.c
@@ -214,12 +214,7 @@ lower_kernel_intrinsics(nir_shader *nir)
             nir_intrinsic_set_range(load, 3 * 4);
             nir_def_init(&load->instr, &load->def, 3, 32);
             nir_builder_instr_insert(&b, &load->instr);
-
-            /* We may need to do a bit-size cast here */
-            nir_def *num_work_groups =
-               nir_u2uN(&b, &load->def, intrin->def.bit_size);
-
-            nir_def_rewrite_uses(&intrin->def, num_work_groups);
+            nir_def_rewrite_uses(&intrin->def, &load->def);
             progress = true;
             break;
          }
diff --git a/src/intel/compiler/brw_nir_lower_cs_intrinsics.c b/src/intel/compiler/brw_nir_lower_cs_intrinsics.c
index fabe16858d2..9088194986f 100644
--- a/src/intel/compiler/brw_nir_lower_cs_intrinsics.c
+++ b/src/intel/compiler/brw_nir_lower_cs_intrinsics.c
@@ -189,19 +189,6 @@ lower_cs_intrinsics_convert_block(struct lower_intrinsics_state *state,
 
       nir_def *sysval;
       switch (intrinsic->intrinsic) {
-      case nir_intrinsic_load_workgroup_size:
-      case nir_intrinsic_load_workgroup_id:
-      case nir_intrinsic_load_num_workgroups:
-         /* Convert this to 32-bit if it's not */
-         if (intrinsic->def.bit_size == 64) {
-            intrinsic->def.bit_size = 32;
-            sysval = nir_u2u64(b, &intrinsic->def);
-            nir_def_rewrite_uses_after(&intrinsic->def,
-                                           sysval,
-                                           sysval->parent_instr);
-         }
-         continue;
-
       case nir_intrinsic_load_local_invocation_index:
       case nir_intrinsic_load_local_invocation_id: {
          if (!local_index && !nir->info.workgroup_size_variable) {