From 30676319c7e3256db465d7d4a1289eea7ee59120 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Marek=20Ol=C5=A1=C3=A1k?= Date: Wed, 11 Jun 2025 17:39:29 -0400 Subject: [PATCH] radeonsi: remove all uses of NIR_PASS_V MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Reviewed-by: Pierre-Eric Pelloux-Prayer Reviewed-by: Timur Kristóf Part-of: --- src/gallium/drivers/radeonsi/si_shader.c | 94 ++++++++++---------- src/gallium/drivers/radeonsi/si_shader_nir.c | 56 ++++++------ 2 files changed, 75 insertions(+), 75 deletions(-) diff --git a/src/gallium/drivers/radeonsi/si_shader.c b/src/gallium/drivers/radeonsi/si_shader.c index 6dbec7d58c6..c8ef138abf2 100644 --- a/src/gallium/drivers/radeonsi/si_shader.c +++ b/src/gallium/drivers/radeonsi/si_shader.c @@ -1045,23 +1045,23 @@ static bool si_lower_io_to_mem(struct si_shader *shader, nir_shader *nir) if (nir->info.stage == MESA_SHADER_VERTEX) { if (key->ge.as_ls) { - NIR_PASS_V(nir, ac_nir_lower_ls_outputs_to_mem, - is_gfx9_mono_tcs ? NULL : si_map_io_driver_location, - sel->screen->info.gfx_level, - key->ge.opt.same_patch_vertices, - is_gfx9_mono_tcs ? next_sel->info.tcs_inputs_via_temp : 0, - is_gfx9_mono_tcs ? next_sel->info.tcs_inputs_via_lds : ~0ull); + NIR_PASS(_, nir, ac_nir_lower_ls_outputs_to_mem, + is_gfx9_mono_tcs ? NULL : si_map_io_driver_location, + sel->screen->info.gfx_level, + key->ge.opt.same_patch_vertices, + is_gfx9_mono_tcs ? next_sel->info.tcs_inputs_via_temp : 0, + is_gfx9_mono_tcs ? next_sel->info.tcs_inputs_via_lds : ~0ull); return true; } else if (key->ge.as_es) { - NIR_PASS_V(nir, ac_nir_lower_es_outputs_to_mem, si_map_io_driver_location, - sel->screen->info.gfx_level, sel->info.esgs_vertex_stride, ~0ULL); + NIR_PASS(_, nir, ac_nir_lower_es_outputs_to_mem, si_map_io_driver_location, + sel->screen->info.gfx_level, sel->info.esgs_vertex_stride, ~0ULL); return true; } } else if (nir->info.stage == MESA_SHADER_TESS_CTRL) { - NIR_PASS_V(nir, ac_nir_lower_hs_inputs_to_mem, - is_gfx9_mono_tcs ? NULL : si_map_io_driver_location, - sel->screen->info.gfx_level, key->ge.opt.same_patch_vertices, - sel->info.tcs_inputs_via_temp, sel->info.tcs_inputs_via_lds); + NIR_PASS(_, nir, ac_nir_lower_hs_inputs_to_mem, + is_gfx9_mono_tcs ? NULL : si_map_io_driver_location, + sel->screen->info.gfx_level, key->ge.opt.same_patch_vertices, + sel->info.tcs_inputs_via_temp, sel->info.tcs_inputs_via_lds); /* Used by hs_emit_write_tess_factors() when monolithic shader. */ if (nir->info.tess._primitive_mode == TESS_PRIMITIVE_UNSPECIFIED) @@ -1074,21 +1074,21 @@ static bool si_lower_io_to_mem(struct si_shader *shader, nir_shader *nir) ac_nir_get_tess_io_info(nir, &tcs_info, ~0ull, ~0, si_map_io_driver_location, false, &tess_io_info); - NIR_PASS_V(nir, ac_nir_lower_hs_outputs_to_mem, &tcs_info, &tess_io_info, si_map_io_driver_location, - sel->screen->info.gfx_level, shader->wave_size); + NIR_PASS(_, nir, ac_nir_lower_hs_outputs_to_mem, &tcs_info, &tess_io_info, si_map_io_driver_location, + sel->screen->info.gfx_level, shader->wave_size); return true; } else if (nir->info.stage == MESA_SHADER_TESS_EVAL) { - NIR_PASS_V(nir, ac_nir_lower_tes_inputs_to_mem, si_map_io_driver_location); + NIR_PASS(_, nir, ac_nir_lower_tes_inputs_to_mem, si_map_io_driver_location); if (key->ge.as_es) { - NIR_PASS_V(nir, ac_nir_lower_es_outputs_to_mem, si_map_io_driver_location, - sel->screen->info.gfx_level, sel->info.esgs_vertex_stride, ~0ULL); + NIR_PASS(_, nir, ac_nir_lower_es_outputs_to_mem, si_map_io_driver_location, + sel->screen->info.gfx_level, sel->info.esgs_vertex_stride, ~0ULL); } return true; } else if (nir->info.stage == MESA_SHADER_GEOMETRY) { - NIR_PASS_V(nir, ac_nir_lower_gs_inputs_to_mem, si_map_io_driver_location, - sel->screen->info.gfx_level, key->ge.mono.u.gs_tri_strip_adj_fix); + NIR_PASS(_, nir, ac_nir_lower_gs_inputs_to_mem, si_map_io_driver_location, + sel->screen->info.gfx_level, key->ge.mono.u.gs_tri_strip_adj_fix); return true; } @@ -1174,8 +1174,8 @@ static void si_lower_ngg(struct si_shader *shader, nir_shader *nir, options.export_primitive_id = key->ge.mono.u.vs_export_prim_id; options.instance_rate_inputs = instance_rate_inputs; - NIR_PASS_V(nir, ac_nir_lower_ngg_nogs, &options, &shader->info.ngg_lds_vertex_size, - &shader->info.ngg_lds_scratch_size); + NIR_PASS(_, nir, ac_nir_lower_ngg_nogs, &options, &shader->info.ngg_lds_vertex_size, + &shader->info.ngg_lds_scratch_size); } else { assert(nir->info.stage == MESA_SHADER_GEOMETRY); @@ -1188,12 +1188,12 @@ static void si_lower_ngg(struct si_shader *shader, nir_shader *nir, if (key->ge.part.gs.es) nir->info.writes_memory |= key->ge.part.gs.es->info.base.writes_memory; - NIR_PASS_V(nir, ac_nir_lower_ngg_gs, &options, &shader->info.ngg_lds_vertex_size, - &shader->info.ngg_lds_scratch_size); + NIR_PASS(_, nir, ac_nir_lower_ngg_gs, &options, &shader->info.ngg_lds_vertex_size, + &shader->info.ngg_lds_scratch_size); } /* may generate some vector output store */ - NIR_PASS_V(nir, nir_lower_io_to_scalar, nir_var_shader_out, NULL, NULL); + NIR_PASS(_, nir, nir_lower_io_to_scalar, nir_var_shader_out, NULL, NULL); } struct nir_shader *si_deserialize_shader(struct si_shader_selector *sel) @@ -1287,8 +1287,8 @@ static void si_assign_param_offsets(nir_shader *nir, struct si_shader *shader, /* This sets DEFAULT_VAL for constant outputs in vs_output_param_offset. */ /* TODO: This doesn't affect GS. */ - NIR_PASS_V(nir, ac_nir_optimize_outputs, false, slot_remap, - temp_info->vs_output_param_offset); + NIR_PASS(_, nir, ac_nir_optimize_outputs, false, slot_remap, + temp_info->vs_output_param_offset); /* Assign the non-constant outputs. */ si_nir_assign_param_offsets(nir, shader, slot_remap, temp_info); @@ -1365,8 +1365,8 @@ static void run_pre_link_optimization_passes(struct si_nir_shader_ctx *ctx) * TODO: The driver uses a linear search to find a shader variant. This * can be really slow if we get too many variants due to uniform inlining. */ - NIR_PASS_V(nir, nir_inline_uniforms, nir->info.num_inlinable_uniforms, - inlined_uniform_values, nir->info.inlinable_uniform_dw_offsets); + NIR_PASS(_, nir, nir_inline_uniforms, nir->info.num_inlinable_uniforms, + inlined_uniform_values, nir->info.inlinable_uniform_dw_offsets); progress = true; } @@ -1459,7 +1459,7 @@ static void run_pre_link_optimization_passes(struct si_nir_shader_ctx *ctx) NIR_PASS(progress, nir, nir_opt_move_to_top, nir_move_to_top_input_loads); /* Remove dead temps before we lower indirect indexing. */ - NIR_PASS_V(nir, nir_remove_dead_variables, nir_var_function_temp, NULL); + NIR_PASS(_, nir, nir_remove_dead_variables, nir_var_function_temp, NULL); /* Lower indirect indexing last. * @@ -1545,15 +1545,15 @@ static void run_late_optimization_and_lowering_passes(struct si_nir_shader_ctx * } else if (nir->info.stage == MESA_SHADER_VERTEX || nir->info.stage == MESA_SHADER_TESS_EVAL) { /* Lower last VGT none-NGG VS/TES shader stage. */ - NIR_PASS_V(nir, ac_nir_lower_legacy_vs, - sel->screen->info.gfx_level, - shader->info.clipdist_mask | shader->info.culldist_mask, - shader->key.ge.mono.write_pos_to_clipvertex, true, - ctx->temp_info.vs_output_param_offset, - shader->info.nr_param_exports, - shader->key.ge.mono.u.vs_export_prim_id, - !shader->info.num_streamout_vec4s, - sel->screen->options.vrs2x2); + NIR_PASS(_, nir, ac_nir_lower_legacy_vs, + sel->screen->info.gfx_level, + shader->info.clipdist_mask | shader->info.culldist_mask, + shader->key.ge.mono.write_pos_to_clipvertex, true, + ctx->temp_info.vs_output_param_offset, + shader->info.nr_param_exports, + shader->key.ge.mono.u.vs_export_prim_id, + !shader->info.num_streamout_vec4s, + sel->screen->options.vrs2x2); } progress = true; } else if (nir->info.stage == MESA_SHADER_GEOMETRY && !key->ge.as_ngg) { @@ -1628,7 +1628,7 @@ static void run_late_optimization_and_lowering_passes(struct si_nir_shader_ctx * if (si_should_clear_lds(sel->screen, nir)) { const unsigned chunk_size = 16; /* max single store size */ const unsigned shared_size = ALIGN(nir->info.shared_size, chunk_size); - NIR_PASS_V(nir, nir_clear_shared_memory, shared_size, chunk_size); + NIR_PASS(_, nir, nir_clear_shared_memory, shared_size, chunk_size); } nir_divergence_analysis(nir); /* required by ac_nir_flag_smem_for_loads */ @@ -1717,7 +1717,7 @@ static void run_late_optimization_and_lowering_passes(struct si_nir_shader_ctx * .buffer_max = ~0, .shared_max = ~0, }; - NIR_PASS_V(nir, nir_opt_offsets, &offset_options); + NIR_PASS(_, nir, nir_opt_offsets, &offset_options); si_nir_late_opts(nir); @@ -1738,7 +1738,7 @@ static void run_late_optimization_and_lowering_passes(struct si_nir_shader_ctx * /* This helps LLVM form VMEM clauses and thus get more GPU cache hits. * 200 is tuned for Viewperf. It should be done last. */ - NIR_PASS_V(nir, nir_group_loads, nir_group_same_resource_only, 200); + NIR_PASS(_, nir, nir_group_loads, nir_group_same_resource_only, 200); } static void get_input_nir(struct si_shader *shader, struct si_nir_shader_ctx *ctx) @@ -1815,7 +1815,7 @@ static void get_nir_shaders(struct si_shader *shader, struct si_linked_shaders * * monolithic PS. */ if (shader->selector->stage == MESA_SHADER_FRAGMENT && shader->is_monolithic) - NIR_PASS_V(linked->consumer.nir, nir_recompute_io_bases, nir_var_shader_in); + NIR_PASS(_, linked->consumer.nir, nir_recompute_io_bases, nir_var_shader_in); for (unsigned i = 0; i < SI_NUM_LINKED_SHADERS; i++) { if (linked->shader[i].nir) { @@ -1860,14 +1860,14 @@ si_nir_generate_gs_copy_shader(struct si_screen *sscreen, si_init_shader_args(shader, &linked.consumer.args, &gs_nir->info); - NIR_PASS_V(nir, si_nir_lower_abi, shader, &linked.consumer.args); - NIR_PASS_V(nir, ac_nir_lower_intrinsics_to_args, sscreen->info.gfx_level, - sscreen->info.has_ls_vgpr_init_bug, AC_HW_VERTEX_SHADER, 64, 64, - &linked.consumer.args.ac); + NIR_PASS(_, nir, si_nir_lower_abi, shader, &linked.consumer.args); + NIR_PASS(_, nir, ac_nir_lower_intrinsics_to_args, sscreen->info.gfx_level, + sscreen->info.has_ls_vgpr_init_bug, AC_HW_VERTEX_SHADER, 64, 64, + &linked.consumer.args.ac); si_nir_opts(gs_selector->screen, nir, false); - NIR_PASS_V(nir, nir_lower_load_const_to_scalar); + NIR_PASS(_, nir, nir_lower_load_const_to_scalar); /* This pass must be last. */ si_get_late_shader_variant_info(shader, &linked.consumer.args, nir); diff --git a/src/gallium/drivers/radeonsi/si_shader_nir.c b/src/gallium/drivers/radeonsi/si_shader_nir.c index fbc20d5c80e..9a11c159356 100644 --- a/src/gallium/drivers/radeonsi/si_shader_nir.c +++ b/src/gallium/drivers/radeonsi/si_shader_nir.c @@ -74,10 +74,10 @@ void si_nir_opts(struct si_screen *sscreen, struct nir_shader *nir, bool has_arr NIR_PASS(progress, nir, nir_opt_dead_cf); if (lower_alu_to_scalar) { - NIR_PASS_V(nir, nir_lower_alu_to_scalar, nir->options->lower_to_scalar_filter, NULL); + NIR_PASS(_, nir, nir_lower_alu_to_scalar, nir->options->lower_to_scalar_filter, NULL); } if (lower_phis_to_scalar) - NIR_PASS_V(nir, nir_lower_phis_to_scalar, NULL, NULL); + NIR_PASS(_, nir, nir_lower_phis_to_scalar, NULL, NULL); progress |= lower_alu_to_scalar | lower_phis_to_scalar; NIR_PASS(progress, nir, nir_opt_cse); @@ -126,13 +126,13 @@ void si_nir_opts(struct si_screen *sscreen, struct nir_shader *nir, bool has_arr } if (nir->info.stage == MESA_SHADER_FRAGMENT) - NIR_PASS_V(nir, nir_opt_move_discards_to_top); + NIR_PASS(_, nir, nir_opt_move_discards_to_top); if (sscreen->info.has_packed_math_16bit) NIR_PASS(progress, nir, nir_opt_vectorize, si_vectorize_callback, NULL); } while (progress); - NIR_PASS_V(nir, nir_lower_var_copies); + NIR_PASS(_, nir, nir_lower_var_copies); } void si_nir_late_opts(nir_shader *nir) @@ -141,18 +141,18 @@ void si_nir_late_opts(nir_shader *nir) while (more_late_algebraic) { more_late_algebraic = false; NIR_PASS(more_late_algebraic, nir, nir_opt_algebraic_late); - NIR_PASS_V(nir, nir_opt_constant_folding); + NIR_PASS(_, nir, nir_opt_constant_folding); /* We should run this after constant folding for stages that support indirect * inputs/outputs. */ if (nir->options->support_indirect_inputs & BITFIELD_BIT(nir->info.stage) || nir->options->support_indirect_outputs & BITFIELD_BIT(nir->info.stage)) - NIR_PASS_V(nir, nir_io_add_const_offset_to_base, nir_var_shader_in | nir_var_shader_out); + NIR_PASS(_, nir, nir_io_add_const_offset_to_base, nir_var_shader_in | nir_var_shader_out); - NIR_PASS_V(nir, nir_copy_prop); - NIR_PASS_V(nir, nir_opt_dce); - NIR_PASS_V(nir, nir_opt_cse); + NIR_PASS(_, nir, nir_copy_prop); + NIR_PASS(_, nir, nir_opt_dce); + NIR_PASS(_, nir, nir_opt_cse); } } @@ -301,30 +301,30 @@ static void si_lower_nir(struct si_screen *sscreen, struct nir_shader *nir) .lower_to_fragment_fetch_amd = sscreen->info.gfx_level < GFX11, .lower_1d = sscreen->info.gfx_level == GFX9, }; - NIR_PASS_V(nir, nir_lower_tex, &lower_tex_options); + NIR_PASS(_, nir, nir_lower_tex, &lower_tex_options); const struct nir_lower_image_options lower_image_options = { .lower_cube_size = true, .lower_to_fragment_mask_load_amd = sscreen->info.gfx_level < GFX11 && !(sscreen->debug_flags & DBG(NO_FMASK)), }; - NIR_PASS_V(nir, nir_lower_image, &lower_image_options); + NIR_PASS(_, nir, nir_lower_image, &lower_image_options); - NIR_PASS_V(nir, si_lower_intrinsics); + NIR_PASS(_, nir, si_lower_intrinsics); - NIR_PASS_V(nir, ac_nir_lower_sin_cos); + NIR_PASS(_, nir, ac_nir_lower_sin_cos); /* Lower load constants to scalar and then clean up the mess */ - NIR_PASS_V(nir, nir_lower_load_const_to_scalar); - NIR_PASS_V(nir, nir_lower_var_copies); - NIR_PASS_V(nir, nir_opt_intrinsics); - NIR_PASS_V(nir, nir_lower_system_values); + NIR_PASS(_, nir, nir_lower_load_const_to_scalar); + NIR_PASS(_, nir, nir_lower_var_copies); + NIR_PASS(_, nir, nir_opt_intrinsics); + NIR_PASS(_, nir, nir_lower_system_values); /* si_nir_kill_outputs and ac_nir_optimize_outputs require outputs to be scalar. */ if (nir->info.stage == MESA_SHADER_VERTEX || nir->info.stage == MESA_SHADER_TESS_EVAL || nir->info.stage == MESA_SHADER_GEOMETRY) - NIR_PASS_V(nir, nir_lower_io_to_scalar, nir_var_shader_out, NULL, NULL); + NIR_PASS(_, nir, nir_lower_io_to_scalar, nir_var_shader_out, NULL, NULL); if (nir->info.stage == MESA_SHADER_GEOMETRY) { unsigned flags = nir_lower_gs_intrinsics_per_stream; @@ -334,7 +334,7 @@ static void si_lower_nir(struct si_screen *sscreen, struct nir_shader *nir) nir_lower_gs_intrinsics_overwrite_incomplete; } - NIR_PASS_V(nir, nir_lower_gs_intrinsics, flags); + NIR_PASS(_, nir, nir_lower_gs_intrinsics, flags); } if (gl_shader_stage_is_compute(nir->info.stage)) { @@ -353,7 +353,7 @@ static void si_lower_nir(struct si_screen *sscreen, struct nir_shader *nir) 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); + NIR_PASS(_, nir, nir_lower_compute_system_values, &options); /* Gfx12 supports this in hw. */ if (sscreen->info.gfx_level < GFX12 && @@ -361,7 +361,7 @@ static void si_lower_nir(struct si_screen *sscreen, struct nir_shader *nir) nir_opt_cse(nir); /* CSE load_local_invocation_id */ memset(&options, 0, sizeof(options)); options.shuffle_local_ids_for_quad_derivatives = true; - NIR_PASS_V(nir, nir_lower_compute_system_values, &options); + NIR_PASS(_, nir, nir_lower_compute_system_values, &options); } } @@ -372,9 +372,9 @@ static void si_lower_nir(struct si_screen *sscreen, struct nir_shader *nir) if (sscreen->info.gfx_level >= GFX9) si_late_optimize_16bit_samplers(sscreen, nir); - NIR_PASS_V(nir, nir_remove_dead_variables, nir_var_function_temp, NULL); + NIR_PASS(_, nir, nir_remove_dead_variables, nir_var_function_temp, NULL); - NIR_PASS_V(nir, nir_lower_fp16_casts, nir_lower_fp16_split_fp64); + NIR_PASS(_, nir, nir_lower_fp16_casts, nir_lower_fp16_split_fp64); } char *si_finalize_nir(struct pipe_screen *screen, struct nir_shader *nir) @@ -387,16 +387,16 @@ char *si_finalize_nir(struct pipe_screen *screen, struct nir_shader *nir) } } else { nir_lower_io_passes(nir, false); - NIR_PASS_V(nir, nir_remove_dead_variables, nir_var_shader_in | nir_var_shader_out, NULL); + NIR_PASS(_, nir, nir_remove_dead_variables, nir_var_shader_in | nir_var_shader_out, NULL); } if (nir->info.stage == MESA_SHADER_FRAGMENT) - NIR_PASS_V(nir, si_nir_lower_color_inputs_to_sysvals); + NIR_PASS(_, nir, si_nir_lower_color_inputs_to_sysvals); - NIR_PASS_V(nir, nir_lower_explicit_io, nir_var_mem_shared, nir_address_format_32bit_offset); + NIR_PASS(_, nir, nir_lower_explicit_io, nir_var_mem_shared, nir_address_format_32bit_offset); /* Remove dead derefs, so that we can remove uniforms. */ - NIR_PASS_V(nir, nir_opt_dce); + NIR_PASS(_, nir, nir_opt_dce); /* Remove uniforms because those should have been lowered to UBOs already. */ nir_foreach_variable_with_modes_safe(var, nir, nir_var_uniform) { @@ -425,7 +425,7 @@ char *si_finalize_nir(struct pipe_screen *screen, struct nir_shader *nir) * nir_opt_large_constants may use op_amul (see nir_build_deref_offset), * or may create unneeded code, so run si_nir_opts if needed. */ - NIR_PASS_V(nir, nir_remove_dead_variables, nir_var_function_temp, NULL); + NIR_PASS(_, nir, nir_remove_dead_variables, nir_var_function_temp, NULL); bool progress = false; NIR_PASS(progress, nir, nir_opt_large_constants, glsl_get_natural_size_align_bytes, 16); if (progress)