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

compiler: Allow derivative_group to be used for all stages in shader_info

Browse Source 
These will now also be used by stages that have workgroups.

Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/30950>
This commit is contained in:
Caio Oliveira
2024-08-30 08:44:23 -07:00
committed by Marge Bot
parent 1eb98bc457
commit 74be809237
16 changed files with 47 additions and 50 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/amd/llvm/ac_nir_to_llvm.c
+2 -2
View File
@@ -1479,7 +1479,7 @@ static LLVMValueRef build_tex_intrinsic(struct ac_nir_context *ctx, const nir_te
case nir_texop_tex:
if (ctx->stage != MESA_SHADER_FRAGMENT &&
(!gl_shader_stage_is_compute(ctx->stage) ||
ctx->info->cs.derivative_group == DERIVATIVE_GROUP_NONE)) {
ctx->info->derivative_group == DERIVATIVE_GROUP_NONE)) {
assert(!args->lod);
args->level_zero = true;
}
@@ -1514,7 +1514,7 @@ static LLVMValueRef build_tex_intrinsic(struct ac_nir_context *ctx, const nir_te
args->attributes = AC_ATTR_INVARIANT_LOAD;
bool cs_derivs =
gl_shader_stage_is_compute(ctx->stage) && ctx->info->cs.derivative_group != DERIVATIVE_GROUP_NONE;
gl_shader_stage_is_compute(ctx->stage) && ctx->info->derivative_group != DERIVATIVE_GROUP_NONE;
if (ctx->stage == MESA_SHADER_FRAGMENT || cs_derivs) {
/* Prevent texture instructions with implicit derivatives from being
* sinked into branches. */
src/compiler/glsl/linker.cpp
+6 -6
View File
@@ -1394,7 +1394,7 @@ link_cs_input_layout_qualifiers(struct gl_shader_program *prog,
gl_prog->info.workgroup_size_variable = false;
gl_prog->info.cs.derivative_group = DERIVATIVE_GROUP_NONE;
gl_prog->info.derivative_group = DERIVATIVE_GROUP_NONE;
/* From the ARB_compute_shader spec, in the section describing local size
* declarations:
@@ -1442,13 +1442,13 @@ link_cs_input_layout_qualifiers(struct gl_shader_program *prog,
enum gl_derivative_group group = shader->info.Comp.DerivativeGroup;
if (group != DERIVATIVE_GROUP_NONE) {
if (gl_prog->info.cs.derivative_group != DERIVATIVE_GROUP_NONE &&
gl_prog->info.cs.derivative_group != group) {
if (gl_prog->info.derivative_group != DERIVATIVE_GROUP_NONE &&
gl_prog->info.derivative_group != group) {
linker_error(prog, "compute shader defined with conflicting "
"derivative groups\n");
return;
}
gl_prog->info.cs.derivative_group = group;
gl_prog->info.derivative_group = group;
}
}
@@ -1463,7 +1463,7 @@ link_cs_input_layout_qualifiers(struct gl_shader_program *prog,
return;
}
if (gl_prog->info.cs.derivative_group == DERIVATIVE_GROUP_QUADS) {
if (gl_prog->info.derivative_group == DERIVATIVE_GROUP_QUADS) {
if (gl_prog->info.workgroup_size[0] % 2 != 0) {
linker_error(prog, "derivative_group_quadsNV must be used with a "
"local group size whose first dimension "
@@ -1476,7 +1476,7 @@ link_cs_input_layout_qualifiers(struct gl_shader_program *prog,
"is a multiple of 2\n");
return;
}
} else if (gl_prog->info.cs.derivative_group == DERIVATIVE_GROUP_LINEAR) {
} else if (gl_prog->info.derivative_group == DERIVATIVE_GROUP_LINEAR) {
if ((gl_prog->info.workgroup_size[0] *
gl_prog->info.workgroup_size[1] *
gl_prog->info.workgroup_size[2]) % 4 != 0) {
src/compiler/nir/nir.c
+1 -1
View File
@@ -2199,7 +2199,7 @@ nir_shader_supports_implicit_lod(nir_shader *shader)
{
return (shader->info.stage == MESA_SHADER_FRAGMENT ||
(shader->info.stage == MESA_SHADER_COMPUTE &&
shader->info.cs.derivative_group != DERIVATIVE_GROUP_NONE));
shader->info.derivative_group != DERIVATIVE_GROUP_NONE));
}
nir_intrinsic_op
src/compiler/nir/nir_builder.h
+1 -1
View File
@@ -2012,7 +2012,7 @@ nir_build_deriv(nir_builder *b, nir_def *x, nir_op alu, nir_intrinsic_op intrin)
* move this to glsl-to-nir.
*/
if (b->shader->info.stage == MESA_SHADER_COMPUTE &&
b->shader->info.cs.derivative_group == DERIVATIVE_GROUP_NONE) {
b->shader->info.derivative_group == DERIVATIVE_GROUP_NONE) {
return nir_imm_zero(b, x->num_components, x->bit_size);
}
src/compiler/nir/nir_lower_system_values.c
+3 -3
View File
@@ -539,7 +539,7 @@ lower_compute_system_value_instr(nir_builder *b,
return lower_id_to_index(b, local_index, local_size, bit_size);
}
if (options && options->shuffle_local_ids_for_quad_derivatives &&
b->shader->info.cs.derivative_group == DERIVATIVE_GROUP_QUADS &&
b->shader->info.derivative_group == DERIVATIVE_GROUP_QUADS &&
_mesa_set_search(state->lower_once_list, instr) == NULL) {
nir_def *ids = nir_load_local_invocation_id(b);
_mesa_set_add(state->lower_once_list, ids->parent_instr);
@@ -803,8 +803,8 @@ nir_lower_compute_system_values(nir_shader *shader,
/* Update this so as not to lower it again. */
if (options && options->shuffle_local_ids_for_quad_derivatives &&
shader->info.cs.derivative_group == DERIVATIVE_GROUP_QUADS)
shader->info.cs.derivative_group = DERIVATIVE_GROUP_LINEAR;
shader->info.derivative_group == DERIVATIVE_GROUP_QUADS)
shader->info.derivative_group = DERIVATIVE_GROUP_LINEAR;
return progress;
}
src/compiler/nir/nir_opt_algebraic.py
+1 -1
View File
@@ -2887,7 +2887,7 @@ for op in ['fadd', 'fmul', 'fmulz', 'iadd', 'imul']:
for op in ['fddx', 'fddx_fine', 'fddx_coarse',
'fddy', 'fddy_fine', 'fddy_coarse']:
optimizations += [
((op, 'a'), 0.0, 'info->stage == MESA_SHADER_COMPUTE && info->cs.derivative_group == DERIVATIVE_GROUP_NONE')
((op, 'a'), 0.0, 'info->stage == MESA_SHADER_COMPUTE && info->derivative_group == DERIVATIVE_GROUP_NONE')
]
# Some optimizations for ir3-specific instructions.
src/compiler/nir/nir_print.c
+2 -2
View File
@@ -1648,7 +1648,7 @@ print_intrinsic_instr(nir_intrinsic_instr *instr, print_state *state)
nir_foreach_variable_with_modes(var, state->shader, var_mode) {
if (!var->name)
continue;
bool match;
if (instr->intrinsic == nir_intrinsic_load_uniform) {
match = var->data.driver_location == nir_intrinsic_base(instr);
@@ -2570,6 +2570,7 @@ print_shader_info(const struct shader_info *info, FILE *fp)
print_nz_bool(fp, "flrp_lowered", info->flrp_lowered);
print_nz_bool(fp, "io_lowered", info->io_lowered);
print_nz_bool(fp, "writes_memory", info->writes_memory);
print_nz_unsigned(fp, "derivative_group", info->derivative_group);
switch (info->stage) {
case MESA_SHADER_VERTEX:
@@ -2649,7 +2650,6 @@ print_shader_info(const struct shader_info *info, FILE *fp)
info->cs.workgroup_size_hint[1],
info->cs.workgroup_size_hint[2]);
print_nz_unsigned(fp, "user_data_components_amd", info->cs.user_data_components_amd);
print_nz_unsigned(fp, "derivative_group", info->cs.derivative_group);
fprintf(fp, "ptr_size: %u\n", info->cs.ptr_size);
break;
src/compiler/shader_info.h
+6 -6
View File
@@ -285,6 +285,12 @@ typedef struct shader_info {
*/
bool use_legacy_math_rules;
/*
* Arrangement of invocations used to calculate derivatives in
* compute/task/mesh shaders. From KHR_compute_shader_derivatives.
*/
enum gl_derivative_group derivative_group:2;
union {
struct {
/* Which inputs are doubles */
@@ -441,12 +447,6 @@ typedef struct shader_info {
uint8_t user_data_components_amd:4;
/*
* Arrangement of invocations used to calculate derivatives in a compute
* shader. From NV_compute_shader_derivatives.
*/
enum gl_derivative_group derivative_group:2;
/*
* If the shader might run with shared mem on top of `shared_size`.
*/
src/compiler/spirv/spirv_to_nir.c
+2 -2
View File
@@ -5274,12 +5274,12 @@ vtn_handle_execution_mode(struct vtn_builder *b, struct vtn_value *entry_point,
case SpvExecutionModeDerivativeGroupQuadsNV:
vtn_assert(b->shader->info.stage == MESA_SHADER_COMPUTE);
b->shader->info.cs.derivative_group = DERIVATIVE_GROUP_QUADS;
b->shader->info.derivative_group = DERIVATIVE_GROUP_QUADS;
break;
case SpvExecutionModeDerivativeGroupLinearNV:
vtn_assert(b->shader->info.stage == MESA_SHADER_COMPUTE);
b->shader->info.cs.derivative_group = DERIVATIVE_GROUP_LINEAR;
b->shader->info.derivative_group = DERIVATIVE_GROUP_LINEAR;
break;
case SpvExecutionModePixelInterlockOrderedEXT:
src/gallium/drivers/radeonsi/si_compute.c
+4 -4
View File
@@ -999,7 +999,7 @@ static void si_emit_dispatch_packets(struct si_context *sctx, const struct pipe_
}
/* Thread tiling within a workgroup. */
switch (sctx->cs_shader_state.program->shader.selector->info.base.cs.derivative_group) {
switch (sctx->cs_shader_state.program->shader.selector->info.base.derivative_group) {
case DERIVATIVE_GROUP_LINEAR:
break;
case DERIVATIVE_GROUP_QUADS:
@@ -1233,10 +1233,10 @@ static void si_launch_grid(struct pipe_context *ctx, const struct pipe_grid_info
NULL);
}
}
if (u_trace_perfetto_active(&sctx->ds.trace_context))
trace_si_begin_compute(&sctx->trace);
if (sctx->bo_list_add_all_compute_resources)
si_compute_resources_add_all_to_bo_list(sctx);
@@ -1310,7 +1310,7 @@ static void si_launch_grid(struct pipe_context *ctx, const struct pipe_grid_info
if (u_trace_perfetto_active(&sctx->ds.trace_context))
trace_si_end_compute(&sctx->trace, info->grid[0], info->grid[1], info->grid[2]);
if (cs_regalloc_hang) {
sctx->flags |= SI_CONTEXT_CS_PARTIAL_FLUSH;
si_mark_atom_dirty(sctx, &sctx->atoms.s.cache_flush);
src/gallium/drivers/radeonsi/si_shader_nir.c
+3 -3
View File
@@ -341,16 +341,16 @@ static void si_lower_nir(struct si_screen *sscreen, struct nir_shader *nir)
* divisible by 2.
*/
options.lower_local_invocation_index =
nir->info.cs.derivative_group == DERIVATIVE_GROUP_QUADS ||
nir->info.derivative_group == DERIVATIVE_GROUP_QUADS ||
(sscreen->info.gfx_level >= GFX12 &&
nir->info.cs.derivative_group == DERIVATIVE_GROUP_NONE &&
nir->info.derivative_group == DERIVATIVE_GROUP_NONE &&
(nir->info.workgroup_size_variable ||
(nir->info.workgroup_size[0] % 2 == 0 && nir->info.workgroup_size[1] % 2 == 0)));
NIR_PASS_V(nir, nir_lower_compute_system_values, &options);
/* Gfx12 supports this in hw. */
if (sscreen->info.gfx_level < GFX12 &&
nir->info.cs.derivative_group == DERIVATIVE_GROUP_QUADS) {
nir->info.derivative_group == DERIVATIVE_GROUP_QUADS) {
nir_opt_cse(nir); /* CSE load_local_invocation_id */
memset(&options, 0, sizeof(options));
options.shuffle_local_ids_for_quad_derivatives = true;
src/gallium/drivers/zink/nir_to_spirv/nir_to_spirv.c
+3 -3
View File
@@ -4763,12 +4763,12 @@ nir_to_spirv(struct nir_shader *s, const struct zink_shader_info *sinfo, const s
spirv_builder_emit_specid(&ctx.builder, ctx.shared_mem_size, ZINK_VARIABLE_SHARED_MEM);
spirv_builder_emit_name(&ctx.builder, ctx.shared_mem_size, "variable_shared_mem");
}
if (s->info.cs.derivative_group) {
if (s->info.derivative_group) {
SpvCapability caps[] = { 0, SpvCapabilityComputeDerivativeGroupQuadsNV, SpvCapabilityComputeDerivativeGroupLinearNV };
SpvExecutionMode modes[] = { 0, SpvExecutionModeDerivativeGroupQuadsNV, SpvExecutionModeDerivativeGroupLinearNV };
spirv_builder_emit_extension(&ctx.builder, "SPV_NV_compute_shader_derivatives");
spirv_builder_emit_cap(&ctx.builder, caps[s->info.cs.derivative_group]);
spirv_builder_emit_exec_mode(&ctx.builder, entry_point, modes[s->info.cs.derivative_group]);
spirv_builder_emit_cap(&ctx.builder, caps[s->info.derivative_group]);
spirv_builder_emit_exec_mode(&ctx.builder, entry_point, modes[s->info.derivative_group]);
ctx.explicit_lod = false;
}
break;
src/intel/compiler/brw_nir_lower_cs_intrinsics.c
+6 -9
View File
@@ -27,7 +27,6 @@
struct lower_intrinsics_state {
nir_shader *nir;
nir_function_impl *impl;
enum gl_derivative_group derivative_group;
bool progress;
bool hw_generated_local_id;
nir_builder builder;
@@ -62,7 +61,7 @@ compute_local_index_id(struct lower_intrinsics_state *state, nir_intrinsic_instr
}
if (state->hw_generated_local_id) {
assert(state->derivative_group != DERIVATIVE_GROUP_QUADS);
assert(nir->info.derivative_group != DERIVATIVE_GROUP_QUADS);
nir_def *local_id_vec = nir_load_local_invocation_id(b);
nir_def *local_id[3] = { nir_channel(b, local_id_vec, 0),
@@ -132,7 +131,7 @@ compute_local_index_id(struct lower_intrinsics_state *state, nir_intrinsic_instr
*/
nir_def *id_x, *id_y, *id_z;
switch (state->derivative_group) {
switch (nir->info.derivative_group) {
case DERIVATIVE_GROUP_NONE:
if (nir->info.num_images == 0 &&
nir->info.num_textures == 0) {
@@ -333,16 +332,14 @@ brw_nir_lower_cs_intrinsics(nir_shader *nir,
struct lower_intrinsics_state state = {
.nir = nir,
.hw_generated_local_id = false,
.derivative_group = gl_shader_stage_is_compute(nir->info.stage) ?
nir->info.cs.derivative_group : DERIVATIVE_GROUP_NONE,
};
/* Constraints from NV_compute_shader_derivatives. */
if (!nir->info.workgroup_size_variable) {
if (state.derivative_group == DERIVATIVE_GROUP_QUADS) {
if (nir->info.derivative_group == DERIVATIVE_GROUP_QUADS) {
assert(nir->info.workgroup_size[0] % 2 == 0);
assert(nir->info.workgroup_size[1] % 2 == 0);
} else if (state.derivative_group == DERIVATIVE_GROUP_LINEAR) {
} else if (nir->info.derivative_group == DERIVATIVE_GROUP_LINEAR) {
ASSERTED unsigned workgroup_size =
nir->info.workgroup_size[0] *
nir->info.workgroup_size[1] *
@@ -353,7 +350,7 @@ brw_nir_lower_cs_intrinsics(nir_shader *nir,
if (devinfo->verx10 >= 125 && prog_data &&
nir->info.stage == MESA_SHADER_COMPUTE &&
state.derivative_group != DERIVATIVE_GROUP_QUADS &&
nir->info.derivative_group != DERIVATIVE_GROUP_QUADS &&
!nir->info.workgroup_size_variable &&
util_is_power_of_two_nonzero(nir->info.workgroup_size[0]) &&
util_is_power_of_two_nonzero(nir->info.workgroup_size[1])) {
@@ -362,7 +359,7 @@ brw_nir_lower_cs_intrinsics(nir_shader *nir,
/* TODO: more heuristics about 1D/SLM access vs. 2D access */
bool linear =
state.derivative_group == DERIVATIVE_GROUP_LINEAR ||
nir->info.derivative_group == DERIVATIVE_GROUP_LINEAR ||
BITSET_TEST(nir->info.system_values_read,
SYSTEM_VALUE_LOCAL_INVOCATION_INDEX) ||
(nir->info.workgroup_size[1] == 1 &&
src/intel/compiler/elk/elk_nir_lower_cs_intrinsics.c
+3 -3
View File
@@ -75,7 +75,7 @@ compute_local_index_id(nir_builder *b,
*/
nir_def *id_x, *id_y, *id_z;
switch (nir->info.cs.derivative_group) {
switch (nir->info.derivative_group) {
case DERIVATIVE_GROUP_NONE:
if (nir->info.num_images == 0 &&
nir->info.num_textures == 0) {
@@ -308,10 +308,10 @@ elk_nir_lower_cs_intrinsics(nir_shader *nir,
/* Constraints from NV_compute_shader_derivatives. */
if (gl_shader_stage_is_compute(nir->info.stage) &&
!nir->info.workgroup_size_variable) {
if (nir->info.cs.derivative_group == DERIVATIVE_GROUP_QUADS) {
if (nir->info.derivative_group == DERIVATIVE_GROUP_QUADS) {
assert(nir->info.workgroup_size[0] % 2 == 0);
assert(nir->info.workgroup_size[1] % 2 == 0);
} else if (nir->info.cs.derivative_group == DERIVATIVE_GROUP_LINEAR) {
} else if (nir->info.derivative_group == DERIVATIVE_GROUP_LINEAR) {
ASSERTED unsigned workgroup_size =
nir->info.workgroup_size[0] *
nir->info.workgroup_size[1] *
src/mesa/main/compute.c
+2 -2
View File
@@ -197,7 +197,7 @@ validate_DispatchComputeGroupSizeARB(struct gl_context *ctx,
* of <group_size_x>, <group_size_y>, and <group_size_z> is not a multiple
* of four."
*/
if (prog->info.cs.derivative_group == DERIVATIVE_GROUP_QUADS &&
if (prog->info.derivative_group == DERIVATIVE_GROUP_QUADS &&
((info->block[0] & 1) || (info->block[1] & 1))) {
_mesa_error(ctx, GL_INVALID_VALUE,
"glDispatchComputeGroupSizeARB(derivative_group_quadsNV "
@@ -206,7 +206,7 @@ validate_DispatchComputeGroupSizeARB(struct gl_context *ctx,
return GL_FALSE;
}
if (prog->info.cs.derivative_group == DERIVATIVE_GROUP_LINEAR &&
if (prog->info.derivative_group == DERIVATIVE_GROUP_LINEAR &&
total_invocations & 3) {
_mesa_error(ctx, GL_INVALID_VALUE,
"glDispatchComputeGroupSizeARB(derivative_group_linearNV "
src/nouveau/compiler/nak_nir.c
+2 -2
View File
@@ -306,7 +306,7 @@ nak_nir_lower_subgroup_id_intrin(nir_builder *b, nir_intrinsic_instr *intrin,
if (nak_nir_workgroup_has_one_subgroup(b->shader)) {
num_subgroups = nir_imm_int(b, 1);
} else {
assert(b->shader->info.cs.derivative_group == DERIVATIVE_GROUP_NONE);
assert(b->shader->info.derivative_group == DERIVATIVE_GROUP_NONE);
nir_def *workgroup_size = nir_load_workgroup_size(b);
workgroup_size =
@@ -327,7 +327,7 @@ nak_nir_lower_subgroup_id_intrin(nir_builder *b, nir_intrinsic_instr *intrin,
if (nak_nir_workgroup_has_one_subgroup(b->shader)) {
subgroup_id = nir_imm_int(b, 0);
} else {
assert(b->shader->info.cs.derivative_group == DERIVATIVE_GROUP_NONE);
assert(b->shader->info.derivative_group == DERIVATIVE_GROUP_NONE);
nir_def *invocation_index = nir_load_local_invocation_index(b);
nir_def *subgroup_size = nir_load_subgroup_size(b);