AlexIndustrial/mesa
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

glsl: handle work group in layout for mesh shader

Browse Source 
Reviewed-by: Ian Romanick <ian.d.romanick@intel.com>
Acked-by: Marek Olšák <marek.olsak@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36405>
This commit is contained in:
Qiang Yu
2025-02-18 15:17:43 +08:00
committed by Marge Bot
parent 611370965f
commit 691601e89e
9 changed files with 114 additions and 36 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/compiler/glsl/ast.h
+5 -5
View File
@@ -1353,14 +1353,14 @@ private:
/**
* AST node representing a decalaration of the input layout for compute
* shaders.
* AST node representing a decalaration of the input layout for compute,
* task and mesh shaders.
*/
class ast_cs_input_layout : public ast_node
class ast_cs_ms_input_layout : public ast_node
{
public:
ast_cs_input_layout(const struct YYLTYPE &locp,
ast_layout_expression *const *local_size)
ast_cs_ms_input_layout(const struct YYLTYPE &locp,
ast_layout_expression *const *local_size)
{
for (int i = 0; i < 3; i++) {
this->local_size[i] = local_size[i];
src/compiler/glsl/ast_to_hir.cpp
+34 -15
View File
@@ -139,7 +139,7 @@ _mesa_ast_to_hir(ir_exec_list *instructions, struct _mesa_glsl_parse_state *stat
state->gs_input_prim_type_specified = false;
state->tcs_output_vertices_specified = false;
state->cs_input_local_size_specified = false;
state->cs_ms_input_local_size_specified = false;
/* Section 4.2 of the GLSL 1.20 specification states:
* "The built-in functions are scoped in a scope outside the global scope
@@ -9002,8 +9002,8 @@ ast_gs_input_layout::hir(ir_exec_list *instructions,
ir_rvalue *
ast_cs_input_layout::hir(ir_exec_list *instructions,
struct _mesa_glsl_parse_state *state)
ast_cs_ms_input_layout::hir(ir_exec_list *instructions,
struct _mesa_glsl_parse_state *state)
{
YYLTYPE loc = this->get_location();
@@ -9018,6 +9018,26 @@ ast_cs_input_layout::hir(ir_exec_list *instructions,
* MAX_COMPUTE_WORK_GROUP_INVOCATIONS, but it seems reasonable to
* report it at compile time as well.
*/
const unsigned *max_work_group_size;
unsigned max_work_group_invocations;
switch (state->stage) {
case MESA_SHADER_COMPUTE:
max_work_group_size = state->consts->MaxComputeWorkGroupSize;
max_work_group_invocations = state->consts->MaxComputeWorkGroupInvocations;
break;
case MESA_SHADER_TASK:
max_work_group_size = state->caps->mesh.max_task_work_group_size;
max_work_group_invocations = state->caps->mesh.max_task_work_group_invocations;
break;
case MESA_SHADER_MESH:
max_work_group_size = state->caps->mesh.max_mesh_work_group_size;
max_work_group_invocations = state->caps->mesh.max_mesh_work_group_invocations;
break;
default:
UNREACHABLE("invalid shader stage");
}
GLuint64 total_invocations = 1;
unsigned qual_local_size[3];
for (int i = 0; i < 3; i++) {
@@ -9035,20 +9055,19 @@ ast_cs_input_layout::hir(ir_exec_list *instructions,
}
ralloc_free(local_size_str);
if (qual_local_size[i] > state->consts->MaxComputeWorkGroupSize[i]) {
if (qual_local_size[i] > max_work_group_size[i]) {
_mesa_glsl_error(&loc, state,
"local_size_%c exceeds MAX_COMPUTE_WORK_GROUP_SIZE"
"local_size_%c exceeds max work group size"
" (%d)", 'x' + i,
state->consts->MaxComputeWorkGroupSize[i]);
max_work_group_size[i]);
break;
}
total_invocations *= qual_local_size[i];
if (total_invocations >
state->consts->MaxComputeWorkGroupInvocations) {
if (total_invocations > max_work_group_invocations) {
_mesa_glsl_error(&loc, state,
"product of local_sizes exceeds "
"MAX_COMPUTE_WORK_GROUP_INVOCATIONS (%d)",
state->consts->MaxComputeWorkGroupInvocations);
"max work group invocations (%d)",
max_work_group_invocations);
break;
}
}
@@ -9056,11 +9075,11 @@ ast_cs_input_layout::hir(ir_exec_list *instructions,
/* If any compute input layout declaration preceded this one, make sure it
* was consistent with this one.
*/
if (state->cs_input_local_size_specified) {
if (state->cs_ms_input_local_size_specified) {
for (int i = 0; i < 3; i++) {
if (state->cs_input_local_size[i] != qual_local_size[i]) {
if (state->cs_ms_input_local_size[i] != qual_local_size[i]) {
_mesa_glsl_error(&loc, state,
"compute shader input layout does not match"
"shader input layout does not match"
" previous declaration");
return NULL;
}
@@ -9081,9 +9100,9 @@ ast_cs_input_layout::hir(ir_exec_list *instructions,
return NULL;
}
state->cs_input_local_size_specified = true;
state->cs_ms_input_local_size_specified = true;
for (int i = 0; i < 3; i++)
state->cs_input_local_size[i] = qual_local_size[i];
state->cs_ms_input_local_size[i] = qual_local_size[i];
/* We may now declare the built-in constant gl_WorkGroupSize (see
* builtin_variable_generator::generate_constants() for why we didn't
src/compiler/glsl/ast_type.cpp
+8 -3
View File
@@ -732,11 +732,16 @@ ast_type_qualifier::validate_in_qualifier(YYLTYPE *loc,
valid_in_mask.flags.q.local_size_variable = 1;
valid_in_mask.flags.q.derivative_group = 1;
break;
case MESA_SHADER_TASK:
case MESA_SHADER_MESH:
valid_in_mask.flags.q.local_size = 7;
break;
default:
r = false;
_mesa_glsl_error(loc, state,
"input layout qualifiers only valid in "
"geometry, tessellation, fragment and compute shaders");
"geometry, tessellation, fragment, task, mesh "
"and compute shaders");
break;
}
@@ -847,8 +852,8 @@ ast_type_qualifier::merge_into_in_qualifier(YYLTYPE *loc,
* into HIR.
*/
if (state->in_qualifier->flags.q.local_size) {
node = new(lin_ctx) ast_cs_input_layout(*loc,
state->in_qualifier->local_size);
node = new(lin_ctx) ast_cs_ms_input_layout(*loc,
state->in_qualifier->local_size);
state->in_qualifier->flags.q.local_size = 0;
for (int i = 0; i < 3; i++)
state->in_qualifier->local_size[i] = NULL;
src/compiler/glsl/builtin_variables.cpp
+1 -1
View File
@@ -941,7 +941,7 @@ builtin_variable_generator::generate_constants()
*
* To prevent the shader from trying to refer to gl_WorkGroupSize before
* the layout declaration, we don't define it here. Intead we define it
* in ast_cs_input_layout::hir().
* in ast_cs_ms_input_layout::hir().
*/
}
src/compiler/glsl/gl_nir_linker.c
+35
View File
@@ -2564,6 +2564,41 @@ link_ms_inout_layout_qualifiers(struct gl_shader_program *prog,
struct gl_shader **shader_list,
unsigned num_shaders)
{
/* handle task and mesh shader in layout */
if (gl_prog->info.stage != MESA_SHADER_MESH &&
gl_prog->info.stage != MESA_SHADER_TASK)
return;
for (int i = 0; i < 3; i++)
gl_prog->nir->info.workgroup_size[i] = 0;
for (unsigned sh = 0; sh < num_shaders; sh++) {
struct gl_shader *shader = shader_list[sh];
if (shader->info.Mesh.LocalSize[0] != 0) {
if (gl_prog->nir->info.workgroup_size[0] != 0) {
for (int i = 0; i < 3; i++) {
if (gl_prog->nir->info.workgroup_size[i] !=
shader->info.Mesh.LocalSize[i]) {
linker_error(prog, "%s shader defined with conflicting local sizes\n",
gl_prog->info.stage == MESA_SHADER_TASK ? "task" : "mesh");
return;
}
}
}
for (int i = 0; i < 3; i++) {
gl_prog->nir->info.workgroup_size[i] = shader->info.Mesh.LocalSize[i];
}
}
}
if (gl_prog->nir->info.workgroup_size[0] == 0) {
linker_error(prog, "%s shader must contain local group size\n",
gl_prog->info.stage == MESA_SHADER_TASK ? "task" : "mesh");
return;
}
/* handle mesh shader out layout */
if (gl_prog->info.stage != MESA_SHADER_MESH)
return;
src/compiler/glsl/glsl_parser.yy
+2 -2
View File
@@ -1867,10 +1867,10 @@ layout_qualifier_id:
for (int i = 0; i < 3; i++) {
if (match_layout_qualifier(local_size_qualifiers[i], $1,
state) == 0) {
if (!state->has_compute_shader()) {
if (!state->has_compute_shader() && !state->EXT_mesh_shader_enable) {
_mesa_glsl_error(& @3, state,
"%s qualifier requires GLSL 4.30 or "
"GLSL ES 3.10 or ARB_compute_shader",
"GLSL ES 3.10 or ARB_compute_shader or EXT_mesh_shader",
local_size_qualifiers[i]);
YYERROR;
} else {
src/compiler/glsl/glsl_parser_extras.cpp
+22 -5
View File
@@ -64,7 +64,7 @@ _mesa_glsl_parse_state::_mesa_glsl_parse_state(struct gl_context *_ctx,
mesa_shader_stage stage,
void *mem_ctx)
: ctx(_ctx), exts(&_ctx->Extensions), consts(&_ctx->Const), caps(&_ctx->screen->caps),
api(_ctx->API), cs_input_local_size_specified(false), cs_input_local_size(),
api(_ctx->API), cs_ms_input_local_size_specified(false), cs_ms_input_local_size(),
switch_state(), warnings_enabled(true)
{
assert(stage < MESA_SHADER_MESH_STAGES);
@@ -1904,13 +1904,20 @@ set_shader_inout_layout(struct gl_shader *shader,
/* Should have been prevented by the parser. */
if (shader->Stage != MESA_SHADER_GEOMETRY &&
shader->Stage != MESA_SHADER_TESS_EVAL &&
shader->Stage != MESA_SHADER_COMPUTE) {
shader->Stage != MESA_SHADER_COMPUTE &&
shader->Stage != MESA_SHADER_TASK &&
shader->Stage != MESA_SHADER_MESH) {
assert(!state->in_qualifier->flags.i);
}
if (shader->Stage != MESA_SHADER_COMPUTE &&
shader->Stage != MESA_SHADER_TASK &&
shader->Stage != MESA_SHADER_MESH) {
assert(!state->cs_ms_input_local_size_specified);
}
if (shader->Stage != MESA_SHADER_COMPUTE) {
/* Should have been prevented by the parser. */
assert(!state->cs_input_local_size_specified);
assert(!state->cs_input_local_size_variable_specified);
assert(state->cs_derivative_group == DERIVATIVE_GROUP_NONE);
}
@@ -2040,9 +2047,9 @@ set_shader_inout_layout(struct gl_shader *shader,
break;
case MESA_SHADER_COMPUTE:
if (state->cs_input_local_size_specified) {
if (state->cs_ms_input_local_size_specified) {
for (int i = 0; i < 3; i++)
shader->info.Comp.LocalSize[i] = state->cs_input_local_size[i];
shader->info.Comp.LocalSize[i] = state->cs_ms_input_local_size[i];
} else {
for (int i = 0; i < 3; i++)
shader->info.Comp.LocalSize[i] = 0;
@@ -2142,6 +2149,16 @@ set_shader_inout_layout(struct gl_shader *shader,
shader->info.Mesh.MaxPrimitives = qual_max_primitives;
}
}
FALLTHROUGH;
case MESA_SHADER_TASK:
if (state->cs_ms_input_local_size_specified) {
for (int i = 0; i < 3; i++)
shader->info.Mesh.LocalSize[i] = state->cs_ms_input_local_size[i];
} else {
for (int i = 0; i < 3; i++)
shader->info.Mesh.LocalSize[i] = 0;
}
break;
default:
src/compiler/glsl/glsl_parser_extras.h
+5 -5
View File
@@ -450,19 +450,19 @@ struct _mesa_glsl_parse_state {
struct ast_type_qualifier *in_qualifier;
/**
* True if a compute shader input local size was specified using a layout
* directive.
* True if a compute, task or mesh shader input local size was specified
* using a layout directive.
*
* Note: this value is computed at ast_to_hir time rather than at parse
* time.
*/
bool cs_input_local_size_specified;
bool cs_ms_input_local_size_specified;
/**
* If cs_input_local_size_specified is true, the local size that was
* If cs_ms_input_local_size_specified is true, the local size that was
* specified. Otherwise ignored.
*/
unsigned cs_input_local_size[3];
unsigned cs_ms_input_local_size[3];
/**
* True if a compute shader input local variable size was specified using
src/mesa/main/shader_types.h
+2
View File
@@ -137,6 +137,8 @@ struct gl_shader_info
* Mesh shader state from EXT_mesh_shader
*/
struct {
unsigned LocalSize[3];
enum mesa_prim OutputType;
GLint MaxVertices;