From f4a65e56287a4cd3492f1e9a3e78db0af90c00c8 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Timur=20Krist=C3=B3f?= Date: Thu, 5 Aug 2021 16:24:35 +0200 Subject: [PATCH] ac/nir/nggc: Only repack arguments that are needed. MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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 Reviewed-by: Daniel Schürmann Part-of: --- src/amd/common/ac_nir.h | 3 +- src/amd/common/ac_nir_lower_ngg.c | 126 ++++++++++++++++++++++++++++-- src/amd/vulkan/radv_shader.c | 3 +- 3 files changed, 123 insertions(+), 9 deletions(-) diff --git a/src/amd/common/ac_nir.h b/src/amd/common/ac_nir.h index 470749949d0..f6d5796bfd1 100644 --- a/src/amd/common/ac_nir.h +++ b/src/amd/common/ac_nir.h @@ -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, diff --git a/src/amd/common/ac_nir_lower_ngg.c b/src/amd/common/ac_nir_lower_ngg.c index c3b525477ff..223d0d71fc6 100644 --- a/src/amd/common/ac_nir_lower_ngg.c +++ b/src/amd/common/ac_nir_lower_ngg.c @@ -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; diff --git a/src/amd/vulkan/radv_shader.c b/src/amd/vulkan/radv_shader.c index 4931ab0f08f..a8f63661576 100644 --- a/src/amd/vulkan/radv_shader.c +++ b/src/amd/vulkan/radv_shader.c @@ -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;