diff --git a/src/mesa/main/compute.c b/src/mesa/main/compute.c
index 9e3f86b6624..fa1cf74d23a 100644
--- a/src/mesa/main/compute.c
+++ b/src/mesa/main/compute.c
@@ -26,6 +26,7 @@
 #include "context.h"
 #include "api_exec_decl.h"
 
+#include "pipe/p_state.h"
 #include "state_tracker/st_cb_compute.h"
 
 static bool
@@ -54,7 +55,7 @@ check_valid_to_compute(struct gl_context *ctx, const char *function)
 }
 
 static bool
-validate_DispatchCompute(struct gl_context *ctx, const GLuint *num_groups)
+validate_DispatchCompute(struct gl_context *ctx, struct pipe_grid_info *info)
 {
    if (!check_valid_to_compute(ctx, "glDispatchCompute"))
       return GL_FALSE;
@@ -78,7 +79,7 @@ validate_DispatchCompute(struct gl_context *ctx, const GLuint *num_groups)
        * Additionally, the OpenGLES 3.1 specification does not contain "or
        * equal to" as an error condition.
        */
-      if (num_groups[i] > ctx->Const.MaxComputeWorkGroupCount[i]) {
+      if (info->grid[i] > ctx->Const.MaxComputeWorkGroupCount[i]) {
          _mesa_error(ctx, GL_INVALID_VALUE,
                      "glDispatchCompute(num_groups_%c)", 'x' + i);
          return GL_FALSE;
@@ -102,8 +103,7 @@ validate_DispatchCompute(struct gl_context *ctx, const GLuint *num_groups)
 
 static bool
 validate_DispatchComputeGroupSizeARB(struct gl_context *ctx,
-                                     const GLuint *num_groups,
-                                     const GLuint *group_size)
+                                     struct pipe_grid_info *info)
 {
    if (!check_valid_to_compute(ctx, "glDispatchComputeGroupSizeARB"))
       return GL_FALSE;
@@ -129,7 +129,7 @@ validate_DispatchComputeGroupSizeARB(struct gl_context *ctx,
        *  num_groups_y and num_groups_z are greater than or equal to the
        *  maximum work group count for the corresponding dimension."
        */
-      if (num_groups[i] > ctx->Const.MaxComputeWorkGroupCount[i]) {
+      if (info->grid[i] > ctx->Const.MaxComputeWorkGroupCount[i]) {
          _mesa_error(ctx, GL_INVALID_VALUE,
                      "glDispatchComputeGroupSizeARB(num_groups_%c)", 'x' + i);
          return GL_FALSE;
@@ -147,8 +147,8 @@ validate_DispatchComputeGroupSizeARB(struct gl_context *ctx,
        * However, the "less than" is a spec bug because they are declared as
        * unsigned integers.
        */
-      if (group_size[i] == 0 ||
-          group_size[i] > ctx->Const.MaxComputeVariableGroupSize[i]) {
+      if (info->block[i] == 0 ||
+          info->block[i] > ctx->Const.MaxComputeVariableGroupSize[i]) {
          _mesa_error(ctx, GL_INVALID_VALUE,
                      "glDispatchComputeGroupSizeARB(group_size_%c)", 'x' + i);
          return GL_FALSE;
@@ -163,19 +163,19 @@ validate_DispatchComputeGroupSizeARB(struct gl_context *ctx,
     *  for compute shaders with variable group size
     *  (MAX_COMPUTE_VARIABLE_GROUP_INVOCATIONS_ARB)."
     */
-   uint64_t total_invocations = group_size[0] * group_size[1];
+   uint64_t total_invocations = info->block[0] * info->block[1];
    if (total_invocations <= UINT32_MAX) {
       /* Only bother multiplying the third value if total still fits in
        * 32-bit, since MaxComputeVariableGroupInvocations is also 32-bit.
        */
-      total_invocations *= group_size[2];
+      total_invocations *= info->block[2];
    }
    if (total_invocations > ctx->Const.MaxComputeVariableGroupInvocations) {
       _mesa_error(ctx, GL_INVALID_VALUE,
                   "glDispatchComputeGroupSizeARB(product of local_sizes "
                   "exceeds MAX_COMPUTE_VARIABLE_GROUP_INVOCATIONS_ARB "
                   "(%u * %u * %u > %u))",
-                  group_size[0], group_size[1], group_size[2],
+                  info->block[0], info->block[1], info->block[2],
                   ctx->Const.MaxComputeVariableGroupInvocations);
       return GL_FALSE;
    }
@@ -194,11 +194,11 @@ validate_DispatchComputeGroupSizeARB(struct gl_context *ctx,
     *  of four."
     */
    if (prog->info.cs.derivative_group == DERIVATIVE_GROUP_QUADS &&
-       ((group_size[0] & 1) || (group_size[1] & 1))) {
+       ((info->block[0] & 1) || (info->block[1] & 1))) {
       _mesa_error(ctx, GL_INVALID_VALUE,
                   "glDispatchComputeGroupSizeARB(derivative_group_quadsNV "
                   "requires group_size_x (%d) and group_size_y (%d) to be "
-                  "divisble by 2)", group_size[0], group_size[1]);
+                  "divisble by 2)", info->block[0], info->block[1]);
       return GL_FALSE;
    }
 
@@ -285,7 +285,7 @@ dispatch_compute(GLuint num_groups_x, GLuint num_groups_y,
                  GLuint num_groups_z, bool no_error)
 {
    GET_CURRENT_CONTEXT(ctx);
-   const GLuint num_groups[3] = { num_groups_x, num_groups_y, num_groups_z };
+   struct pipe_grid_info info = { 0 };
 
    FLUSH_VERTICES(ctx, 0, 0);
 
@@ -293,13 +293,23 @@ dispatch_compute(GLuint num_groups_x, GLuint num_groups_y,
       _mesa_debug(ctx, "glDispatchCompute(%d, %d, %d)\n",
                   num_groups_x, num_groups_y, num_groups_z);
 
-   if (!no_error && !validate_DispatchCompute(ctx, num_groups))
+   info.grid[0] = num_groups_x;
+   info.grid[1] = num_groups_y;
+   info.grid[2] = num_groups_z;
+
+   if (!no_error && !validate_DispatchCompute(ctx, &info))
       return;
 
    if (num_groups_x == 0u || num_groups_y == 0u || num_groups_z == 0u)
        return;
 
-   st_dispatch_compute(ctx, num_groups);
+   struct gl_program *prog =
+      ctx->_Shader->CurrentProgram[MESA_SHADER_COMPUTE];
+   info.block[0] = prog->info.workgroup_size[0];
+   info.block[1] = prog->info.workgroup_size[1];
+   info.block[2] = prog->info.workgroup_size[2];
+
+   st_dispatch_compute(ctx, &info);
 
    if (MESA_DEBUG_FLAGS & DEBUG_ALWAYS_FLUSH)
       _mesa_flush(ctx);
@@ -333,7 +343,17 @@ dispatch_compute_indirect(GLintptr indirect, bool no_error)
    if (!no_error && !valid_dispatch_indirect(ctx, indirect))
       return;
 
-   st_dispatch_compute_indirect(ctx, indirect);
+   struct pipe_grid_info info = { 0 };
+   info.indirect_offset = indirect;
+   info.indirect = ctx->DispatchIndirectBuffer->buffer;
+
+   struct gl_program *prog =
+      ctx->_Shader->CurrentProgram[MESA_SHADER_COMPUTE];
+   info.block[0] = prog->info.workgroup_size[0];
+   info.block[1] = prog->info.workgroup_size[1];
+   info.block[2] = prog->info.workgroup_size[2];
+
+   st_dispatch_compute(ctx, &info);
 
    if (MESA_DEBUG_FLAGS & DEBUG_ALWAYS_FLUSH)
       _mesa_flush(ctx);
@@ -358,9 +378,6 @@ dispatch_compute_group_size(GLuint num_groups_x, GLuint num_groups_y,
                             bool no_error)
 {
    GET_CURRENT_CONTEXT(ctx);
-   const GLuint num_groups[3] = { num_groups_x, num_groups_y, num_groups_z };
-   const GLuint group_size[3] = { group_size_x, group_size_y, group_size_z };
-
    FLUSH_VERTICES(ctx, 0, 0);
 
    if (MESA_VERBOSE & VERBOSE_API)
@@ -369,14 +386,23 @@ dispatch_compute_group_size(GLuint num_groups_x, GLuint num_groups_y,
                   num_groups_x, num_groups_y, num_groups_z,
                   group_size_x, group_size_y, group_size_z);
 
+   struct pipe_grid_info info = { 0 };
+   info.grid[0] = num_groups_x;
+   info.grid[1] = num_groups_y;
+   info.grid[2] = num_groups_z;
+
+   info.block[0] = group_size_x;
+   info.block[1] = group_size_y;
+   info.block[2] = group_size_z;
+
    if (!no_error &&
-       !validate_DispatchComputeGroupSizeARB(ctx, num_groups, group_size))
+       !validate_DispatchComputeGroupSizeARB(ctx, &info))
       return;
 
    if (num_groups_x == 0u || num_groups_y == 0u || num_groups_z == 0u)
        return;
 
-   st_dispatch_compute_group_size(ctx, num_groups, group_size);
+   st_dispatch_compute(ctx, &info);
 
    if (MESA_DEBUG_FLAGS & DEBUG_ALWAYS_FLUSH)
       _mesa_flush(ctx);
diff --git a/src/mesa/state_tracker/st_cb_compute.c b/src/mesa/state_tracker/st_cb_compute.c
index ea6d893f9ee..79040343397 100644
--- a/src/mesa/state_tracker/st_cb_compute.c
+++ b/src/mesa/state_tracker/st_cb_compute.c
@@ -34,17 +34,11 @@
 
 #include "pipe/p_context.h"
 
-static void st_dispatch_compute_common(struct gl_context *ctx,
-                                       const GLuint *num_groups,
-                                       const GLuint *group_size,
-                                       struct pipe_resource *indirect,
-                                       GLintptr indirect_offset)
+void st_dispatch_compute(struct gl_context *ctx,
+                         struct pipe_grid_info *info)
 {
-   struct gl_program *prog =
-      ctx->_Shader->CurrentProgram[MESA_SHADER_COMPUTE];
    struct st_context *st = st_context(ctx);
    struct pipe_context *pipe = st->pipe;
-   struct pipe_grid_info info = { 0 };
 
    st_flush_bitmap_cache(st);
    st_invalidate_readpix_cache(st);
@@ -57,38 +51,5 @@ static void st_dispatch_compute_common(struct gl_context *ctx,
        st->compute_shader_may_be_dirty)
       st_validate_state(st, ST_PIPELINE_COMPUTE);
 
-   for (unsigned i = 0; i < 3; i++) {
-      info.block[i] = group_size ? group_size[i] : prog->info.workgroup_size[i];
-      info.grid[i]  = num_groups ? num_groups[i] : 0;
-   }
-
-   if (indirect) {
-      info.indirect = indirect;
-      info.indirect_offset = indirect_offset;
-   }
-
-   pipe->launch_grid(pipe, &info);
+   pipe->launch_grid(pipe, info);
 }
-
-void st_dispatch_compute(struct gl_context *ctx,
-                         const GLuint *num_groups)
-{
-   st_dispatch_compute_common(ctx, num_groups, NULL, NULL, 0);
-}
-
-void st_dispatch_compute_indirect(struct gl_context *ctx,
-                                  GLintptr indirect_offset)
-{
-   struct gl_buffer_object *indirect_buffer = ctx->DispatchIndirectBuffer;
-   struct pipe_resource *indirect = indirect_buffer->buffer;
-
-   st_dispatch_compute_common(ctx, NULL, NULL, indirect, indirect_offset);
-}
-
-void st_dispatch_compute_group_size(struct gl_context *ctx,
-                                    const GLuint *num_groups,
-                                    const GLuint *group_size)
-{
-   st_dispatch_compute_common(ctx, num_groups, group_size, NULL, 0);
-}
-
diff --git a/src/mesa/state_tracker/st_cb_compute.h b/src/mesa/state_tracker/st_cb_compute.h
index 1f686b5ce40..825ac9e26cd 100644
--- a/src/mesa/state_tracker/st_cb_compute.h
+++ b/src/mesa/state_tracker/st_cb_compute.h
@@ -28,12 +28,9 @@
 #ifndef ST_CB_COMPUTE_H
 #define ST_CB_COMPUTE_H
 
+struct pipe_grid_info;
+
 void st_dispatch_compute(struct gl_context *ctx,
-                         const GLuint *num_groups);
-void st_dispatch_compute_indirect(struct gl_context *ctx,
-                                  GLintptr indirect_offset);
-void st_dispatch_compute_group_size(struct gl_context *ctx,
-                                    const GLuint *num_groups,
-                                    const GLuint *group_size);
+                         struct pipe_grid_info *grid_info);
 
 #endif /* ST_CB_COMPUTE_H */