diff --git a/src/broadcom/compiler/nir_to_vir.c b/src/broadcom/compiler/nir_to_vir.c
index b6269f7e85d..0b2f7c5455e 100644
--- a/src/broadcom/compiler/nir_to_vir.c
+++ b/src/broadcom/compiler/nir_to_vir.c
@@ -3668,6 +3668,18 @@ ntq_emit_intrinsic(struct v3d_compile *c, nir_intrinsic_instr *instr)
                 break;
         }
 
+        case nir_intrinsic_load_workgroup_size: {
+                struct qreg x = vir_uniform(c, QUNIFORM_WORK_GROUP_SIZE, 0);
+                ntq_store_def(c, &instr->def, 0, x);
+
+                struct qreg y = vir_uniform(c, QUNIFORM_WORK_GROUP_SIZE, 1);
+                ntq_store_def(c, &instr->def, 1, y);
+
+                struct qreg z = vir_uniform(c, QUNIFORM_WORK_GROUP_SIZE, 2);
+                ntq_store_def(c, &instr->def, 2, z);
+                break;
+        }
+
         case nir_intrinsic_load_local_invocation_index:
                 ntq_store_def(c, &instr->def, 0,
                               emit_load_local_invocation_index(c));
diff --git a/src/broadcom/compiler/v3d_compiler.h b/src/broadcom/compiler/v3d_compiler.h
index 8b5805652ba..ae8ede54917 100644
--- a/src/broadcom/compiler/v3d_compiler.h
+++ b/src/broadcom/compiler/v3d_compiler.h
@@ -311,6 +311,9 @@ enum quniform_contents {
          */
         QUNIFORM_WORK_GROUP_BASE,
 
+        /* Workgroup size for variable workgroup support */
+        QUNIFORM_WORK_GROUP_SIZE,
+
         /**
          * Returns the the offset of the scratch buffer for register spilling.
          */