Mesa (main): mesa/st: refactor compute dispatch to fill grid info earlier.
GitLab Mirror
gitlab-mirror at kemper.freedesktop.org
Mon Dec 20 04:34:29 UTC 2021
Module: Mesa
Branch: main
Commit: 20de14c57ecb44dea3194dac7fe059e2511a2468
URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=20de14c57ecb44dea3194dac7fe059e2511a2468
Author: Dave Airlie <airlied at redhat.com>
Date: Thu Dec 9 10:56:49 2021 +1000
mesa/st: refactor compute dispatch to fill grid info earlier.
This fills the grid info earlier and uses info in validation
Reviewed-by: Emma Anholt <emma at anholt.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/14256>
---
src/mesa/main/compute.c | 68 +++++++++++++++++++++++-----------
src/mesa/state_tracker/st_cb_compute.c | 45 ++--------------------
src/mesa/state_tracker/st_cb_compute.h | 9 ++---
3 files changed, 53 insertions(+), 69 deletions(-)
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);
-}
-
-void st_dispatch_compute(struct gl_context *ctx,
- const GLuint *num_groups)
-{
- st_dispatch_compute_common(ctx, num_groups, NULL, NULL, 0);
+ pipe->launch_grid(pipe, info);
}
-
-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 */
More information about the mesa-commit
mailing list