fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2025-12-26 17:10:11 +01:00
Code Issues Projects Releases Packages Wiki Activity Actions 1

mesa/st: refactor compute dispatch to fill grid info earlier.

Browse source 
This fills the grid info earlier and uses info in validation

Reviewed-by: Emma Anholt <emma@anholt.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/14256>
This commit is contained in:
Dave Airlie 2021-12-09 10:56:49 +10:00
parent e9cc1633a2
commit 20de14c57e
3 changed files with 53 additions and 69 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

68
src/mesa/main/compute.c
View file

@ -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);

45
src/mesa/state_tracker/st_cb_compute.c
View file

@ -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);
}

9
src/mesa/state_tracker/st_cb_compute.h
View file

@ -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 */