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

ac/nir/nggc: Only repack arguments that are needed.

Browse Source 
Don't repack everything, only what is actually used.
The goal of this commit is primarily to remove unnecessary
LDS stores and loads. In addition to that, it also gets rid of
a few VALU instructions and reduces VGPR use.

Fossil DB stats on Sienna Cichlid with NGGC on:

Totals from 6951 (5.40% of 128647) affected shaders:
VGPRs: 206056 -> 205360 (-0.34%); split: -0.79%, +0.45%
CodeSize: 12344568 -> 12269312 (-0.61%); split: -0.62%, +0.01%
MaxWaves: 211206 -> 212196 (+0.47%)
Instrs: 2319459 -> 2308483 (-0.47%); split: -0.50%, +0.03%
Latency: 7220829 -> 7164721 (-0.78%); split: -1.21%, +0.43%
InvThroughput: 1051450 -> 1049191 (-0.21%); split: -0.36%, +0.15%
VClause: 25794 -> 25445 (-1.35%); split: -1.97%, +0.61%
SClause: 39192 -> 39277 (+0.22%); split: -0.21%, +0.43%
Copies: 315756 -> 313404 (-0.74%); split: -1.17%, +0.42%
Branches: 127878 -> 127879 (+0.00%); split: -0.00%, +0.00%
PreVGPRs: 168029 -> 160162 (-4.68%)

Signed-off-by: Timur Kristóf <timur.kristof@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/12246>
This commit is contained in:
Timur Kristóf
2021-08-05 16:24:35 +02:00
committed by Marge Bot
parent 02bba6aab5
commit f4a65e5628
3 changed files with 123 additions and 9 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/amd/common/ac_nir.h
+2 -1
View File
@@ -110,7 +110,8 @@ ac_nir_lower_ngg_nogs(nir_shader *shader,
bool consider_culling,
bool consider_passthrough,
bool export_prim_id,
bool provoking_vtx_last);
bool provoking_vtx_last,
uint32_t instance_rate_inputs);
void
ac_nir_lower_ngg_gs(nir_shader *shader,
src/amd/common/ac_nir_lower_ngg.c
+119 -7
View File
@@ -60,6 +60,10 @@ typedef struct
uint64_t inputs_needed_by_pos;
uint64_t inputs_needed_by_others;
uint32_t instance_rate_inputs;
nir_instr *compact_arg_stores[4];
nir_intrinsic_instr *overwrite_args;
} lower_ngg_nogs_state;
typedef struct
@@ -526,6 +530,105 @@ remove_extra_pos_outputs(nir_shader *shader, lower_ngg_nogs_state *nogs_state)
nir_metadata_block_index | nir_metadata_dominance, &s);
}
static bool
remove_compacted_arg(lower_ngg_nogs_state *state, nir_builder *b, unsigned idx)
{
nir_instr *store_instr = state->compact_arg_stores[idx];
if (!store_instr)
return false;
/* Simply remove the store. */
nir_instr_remove(store_instr);
/* Find the intrinsic that overwrites the shader arguments,
* and change its corresponding source.
* This will cause NIR's DCE to recognize the load and its phis as dead.
*/
b->cursor = nir_before_instr(&state->overwrite_args->instr);
nir_ssa_def *undef_arg = nir_ssa_undef(b, 1, 32);
nir_ssa_def_rewrite_uses(state->overwrite_args->src[idx].ssa, undef_arg);
state->compact_arg_stores[idx] = NULL;
return true;
}
static bool
cleanup_culling_shader_after_dce(nir_shader *shader,
nir_function_impl *function_impl,
lower_ngg_nogs_state *state)
{
bool uses_vs_vertex_id = false;
bool uses_vs_instance_id = false;
bool uses_tes_u = false;
bool uses_tes_v = false;
bool uses_tes_rel_patch_id = false;
bool uses_tes_patch_id = false;
bool progress = false;
nir_builder b;
nir_builder_init(&b, function_impl);
nir_foreach_block_reverse_safe(block, function_impl) {
nir_foreach_instr_reverse_safe(instr, block) {
if (instr->type != nir_instr_type_intrinsic)
continue;
nir_intrinsic_instr *intrin = nir_instr_as_intrinsic(instr);
switch (intrin->intrinsic) {
case nir_intrinsic_alloc_vertices_and_primitives_amd:
goto cleanup_culling_shader_after_dce_done;
case nir_intrinsic_load_vertex_id:
case nir_intrinsic_load_vertex_id_zero_base:
uses_vs_vertex_id = true;
break;
case nir_intrinsic_load_instance_id:
uses_vs_instance_id = true;
break;
case nir_intrinsic_load_input:
if (state->instance_rate_inputs &
(1 << (nir_intrinsic_base(intrin) - VERT_ATTRIB_GENERIC0)))
uses_vs_instance_id = true;
else
uses_vs_vertex_id = true;
break;
case nir_intrinsic_load_tess_coord:
uses_tes_u = uses_tes_v = true;
break;
case nir_intrinsic_load_tess_rel_patch_id_amd:
uses_tes_rel_patch_id = true;
break;
case nir_intrinsic_load_primitive_id:
if (shader->info.stage == MESA_SHADER_TESS_EVAL)
uses_tes_patch_id = true;
break;
default:
break;
}
}
}
cleanup_culling_shader_after_dce_done:
if (shader->info.stage == MESA_SHADER_VERTEX) {
if (!uses_vs_vertex_id)
progress |= remove_compacted_arg(state, &b, 0);
if (!uses_vs_instance_id)
progress |= remove_compacted_arg(state, &b, 1);
} else if (shader->info.stage == MESA_SHADER_TESS_EVAL) {
if (!uses_tes_u)
progress |= remove_compacted_arg(state, &b, 0);
if (!uses_tes_v)
progress |= remove_compacted_arg(state, &b, 1);
if (!uses_tes_rel_patch_id)
progress |= remove_compacted_arg(state, &b, 2);
if (!uses_tes_patch_id)
progress |= remove_compacted_arg(state, &b, 3);
}
return progress;
}
/**
* Perform vertex compaction after culling.
*
@@ -573,7 +676,9 @@ compact_vertices_after_culling(nir_builder *b,
/* Store the current thread's repackable arguments to the exporter thread's LDS space */
for (unsigned i = 0; i < max_exported_args; ++i) {
nir_ssa_def *arg_val = nir_load_var(b, repacked_arg_vars[i]);
nir_build_store_shared(b, arg_val, exporter_addr, .base = lds_es_arg_0 + 4u * i, .align_mul = 4u, .write_mask = 0x1u);
nir_intrinsic_instr *store = nir_build_store_shared(b, arg_val, exporter_addr, .base = lds_es_arg_0 + 4u * i, .align_mul = 4u, .write_mask = 0x1u);
nogs_state->compact_arg_stores[i] = &store->instr;
}
}
nir_pop_if(b, if_es_accepted);
@@ -1076,12 +1181,14 @@ add_deferred_attribute_culling(nir_builder *b, nir_cf_list *original_extracted_c
*/
if (b->shader->info.stage == MESA_SHADER_VERTEX)
nir_build_overwrite_vs_arguments_amd(b,
nir_load_var(b, repacked_arg_vars[0]), nir_load_var(b, repacked_arg_vars[1]));
nogs_state->overwrite_args =
nir_build_overwrite_vs_arguments_amd(b,
nir_load_var(b, repacked_arg_vars[0]), nir_load_var(b, repacked_arg_vars[1]));
else if (b->shader->info.stage == MESA_SHADER_TESS_EVAL)
nir_build_overwrite_tes_arguments_amd(b,
nir_load_var(b, repacked_arg_vars[0]), nir_load_var(b, repacked_arg_vars[1]),
nir_load_var(b, repacked_arg_vars[2]), nir_load_var(b, repacked_arg_vars[3]));
nogs_state->overwrite_args =
nir_build_overwrite_tes_arguments_amd(b,
nir_load_var(b, repacked_arg_vars[0]), nir_load_var(b, repacked_arg_vars[1]),
nir_load_var(b, repacked_arg_vars[2]), nir_load_var(b, repacked_arg_vars[3]));
else
unreachable("Should be VS or TES.");
}
@@ -1116,7 +1223,8 @@ ac_nir_lower_ngg_nogs(nir_shader *shader,
bool consider_culling,
bool consider_passthrough,
bool export_prim_id,
bool provoking_vtx_last)
bool provoking_vtx_last,
uint32_t instance_rate_inputs)
{
nir_function_impl *impl = nir_shader_get_entrypoint(shader);
assert(impl);
@@ -1145,6 +1253,7 @@ ac_nir_lower_ngg_nogs(nir_shader *shader,
.max_num_waves = DIV_ROUND_UP(max_workgroup_size, wave_size),
.max_es_num_vertices = max_num_es_vertices,
.wave_size = wave_size,
.instance_rate_inputs = instance_rate_inputs,
};
/* We need LDS space when VS needs to export the primitive ID. */
@@ -1260,6 +1369,9 @@ ac_nir_lower_ngg_nogs(nir_shader *shader,
NIR_PASS(progress, shader, nir_opt_undef);
NIR_PASS(progress, shader, nir_opt_dce);
NIR_PASS(progress, shader, nir_opt_dead_cf);
if (can_cull)
progress |= cleanup_culling_shader_after_dce(shader, b->impl, &state);
} while (progress);
shader->info.shared_size = state.total_lds_bytes;
src/amd/vulkan/radv_shader.c
+2 -1
View File
@@ -972,7 +972,8 @@ void radv_lower_ngg(struct radv_device *device, struct nir_shader *nir,
consider_culling,
key->vs_common_out.as_ngg_passthrough,
key->vs_common_out.export_prim_id,
key->vs.provoking_vtx_last);
key->vs.provoking_vtx_last,
key->vs.instance_rate_inputs);
info->has_ngg_culling = out_conf.can_cull;
info->has_ngg_early_prim_export = out_conf.early_prim_export;