diff --git a/src/amd/common/ac_nir.c b/src/amd/common/ac_nir.c index c2f2d66ba54..6c8db43c516 100644 --- a/src/amd/common/ac_nir.c +++ b/src/amd/common/ac_nir.c @@ -74,6 +74,6 @@ ac_nir_lower_indirect_derefs(nir_shader *shader, */ indirect_mask |= nir_var_function_temp; - progress |= nir_lower_indirect_derefs(shader, indirect_mask, UINT32_MAX); + NIR_PASS(progress, shader, nir_lower_indirect_derefs, indirect_mask, UINT32_MAX); return progress; } diff --git a/src/amd/vulkan/radv_pipeline.c b/src/amd/vulkan/radv_pipeline.c index 1fb2fdc2116..71760293c09 100644 --- a/src/amd/vulkan/radv_pipeline.c +++ b/src/amd/vulkan/radv_pipeline.c @@ -2840,18 +2840,20 @@ radv_link_shaders(struct radv_pipeline *pipeline, if (ordered_shaders[i]->info.stage != last) mask = mask | nir_var_shader_out; - if (nir_lower_io_to_scalar_early(ordered_shaders[i], mask)) { + bool progress = false; + NIR_PASS(progress, ordered_shaders[i], nir_lower_io_to_scalar_early, mask); + if (progress) { /* Optimize the new vector code and then remove dead vars */ - nir_copy_prop(ordered_shaders[i]); - nir_opt_shrink_vectors(ordered_shaders[i]); + NIR_PASS(_, ordered_shaders[i], nir_copy_prop); + NIR_PASS(_, ordered_shaders[i], nir_opt_shrink_vectors); if (ordered_shaders[i]->info.stage != last) { /* Optimize swizzled movs of load_const for * nir_link_opt_varyings's constant propagation */ - nir_opt_constant_folding(ordered_shaders[i]); + NIR_PASS(_, ordered_shaders[i], nir_opt_constant_folding); /* For nir_link_opt_varyings's duplicate input opt */ - nir_opt_cse(ordered_shaders[i]); + NIR_PASS(_, ordered_shaders[i], nir_opt_cse); } /* Run copy-propagation to help remove dead @@ -2863,12 +2865,11 @@ radv_link_shaders(struct radv_pipeline *pipeline, * not have worked because the outputs were vector. */ if (ordered_shaders[i]->info.stage == MESA_SHADER_TESS_CTRL) - nir_opt_copy_prop_vars(ordered_shaders[i]); + NIR_PASS(_, ordered_shaders[i], nir_opt_copy_prop_vars); - nir_opt_dce(ordered_shaders[i]); - nir_remove_dead_variables( - ordered_shaders[i], nir_var_function_temp | nir_var_shader_in | nir_var_shader_out, - NULL); + NIR_PASS(_, ordered_shaders[i], nir_opt_dce); + NIR_PASS(_, ordered_shaders[i], nir_remove_dead_variables, + nir_var_function_temp | nir_var_shader_in | nir_var_shader_out, NULL); } } } @@ -2903,9 +2904,10 @@ radv_link_shaders(struct radv_pipeline *pipeline, } } if (fixup_derefs) { - nir_fixup_deref_modes(ordered_shaders[i]); - nir_remove_dead_variables(ordered_shaders[i], nir_var_shader_temp, NULL); - nir_opt_dce(ordered_shaders[i]); + NIR_PASS_V(ordered_shaders[i], nir_fixup_deref_modes); + NIR_PASS(_, ordered_shaders[i], nir_remove_dead_variables, nir_var_shader_temp, + NULL); + NIR_PASS(_, ordered_shaders[i], nir_opt_dce); } continue; } @@ -2935,9 +2937,9 @@ radv_link_shaders(struct radv_pipeline *pipeline, psiz_var->data.mode = nir_var_shader_temp; info->outputs_written &= ~VARYING_BIT_PSIZ; - nir_fixup_deref_modes(ordered_shaders[i]); - nir_remove_dead_variables(ordered_shaders[i], nir_var_shader_temp, NULL); - nir_opt_dce(ordered_shaders[i]); + NIR_PASS_V(ordered_shaders[i], nir_fixup_deref_modes); + NIR_PASS(_, ordered_shaders[i], nir_remove_dead_variables, nir_var_shader_temp, NULL); + NIR_PASS(_, ordered_shaders[i], nir_opt_dce); } } } @@ -2946,7 +2948,7 @@ radv_link_shaders(struct radv_pipeline *pipeline, if (stages[MESA_SHADER_FRAGMENT].nir && (stages[MESA_SHADER_FRAGMENT].nir->info.inputs_read & VARYING_BIT_VIEWPORT) && !(stages[pipeline->graphics.last_vgt_api_stage].nir->info.outputs_written & VARYING_BIT_VIEWPORT)) { - radv_lower_viewport_to_zero(stages[MESA_SHADER_FRAGMENT].nir); + NIR_PASS(_, stages[MESA_SHADER_FRAGMENT].nir, radv_lower_viewport_to_zero); } /* Export the layer in the last VGT stage if multiview is used. */ @@ -2954,18 +2956,18 @@ radv_link_shaders(struct radv_pipeline *pipeline, !(stages[pipeline->graphics.last_vgt_api_stage].nir->info.outputs_written & VARYING_BIT_LAYER)) { nir_shader *last_vgt_shader = stages[pipeline->graphics.last_vgt_api_stage].nir; - radv_lower_multiview(last_vgt_shader); + NIR_PASS(_, last_vgt_shader, radv_lower_multiview); } for (int i = 1; !optimize_conservatively && (i < shader_count); ++i) { if (nir_link_opt_varyings(ordered_shaders[i], ordered_shaders[i - 1])) { - nir_opt_constant_folding(ordered_shaders[i - 1]); - nir_opt_algebraic(ordered_shaders[i - 1]); - nir_opt_dce(ordered_shaders[i - 1]); + NIR_PASS(_, ordered_shaders[i - 1], nir_opt_constant_folding); + NIR_PASS(_, ordered_shaders[i - 1], nir_opt_algebraic); + NIR_PASS(_, ordered_shaders[i - 1], nir_opt_dce); } - nir_remove_dead_variables(ordered_shaders[i], nir_var_shader_out, NULL); - nir_remove_dead_variables(ordered_shaders[i - 1], nir_var_shader_in, NULL); + NIR_PASS(_, ordered_shaders[i], nir_remove_dead_variables, nir_var_shader_out, NULL); + NIR_PASS(_, ordered_shaders[i - 1], nir_remove_dead_variables, nir_var_shader_in, NULL); bool progress = nir_remove_unused_varyings(ordered_shaders[i], ordered_shaders[i - 1]); @@ -2979,27 +2981,31 @@ radv_link_shaders(struct radv_pipeline *pipeline, ordered_shaders[i]->info.stage == MESA_SHADER_MESH || (ordered_shaders[i]->info.stage == MESA_SHADER_VERTEX && has_geom_tess) || (ordered_shaders[i]->info.stage == MESA_SHADER_TESS_EVAL && merged_gs)) { - nir_lower_io_to_vector(ordered_shaders[i], nir_var_shader_out); + NIR_PASS(_, ordered_shaders[i], nir_lower_io_to_vector, nir_var_shader_out); if (ordered_shaders[i]->info.stage == MESA_SHADER_TESS_CTRL) - nir_vectorize_tess_levels(ordered_shaders[i]); - nir_opt_combine_stores(ordered_shaders[i], nir_var_shader_out); + NIR_PASS(_, ordered_shaders[i], nir_vectorize_tess_levels); + NIR_PASS(_, ordered_shaders[i], nir_opt_combine_stores, nir_var_shader_out); } if (ordered_shaders[i - 1]->info.stage == MESA_SHADER_GEOMETRY || ordered_shaders[i - 1]->info.stage == MESA_SHADER_TESS_CTRL || ordered_shaders[i - 1]->info.stage == MESA_SHADER_TESS_EVAL) { - nir_lower_io_to_vector(ordered_shaders[i - 1], nir_var_shader_in); + NIR_PASS(_, ordered_shaders[i - 1], nir_lower_io_to_vector, nir_var_shader_in); } if (progress) { - if (nir_lower_global_vars_to_local(ordered_shaders[i])) { + progress = false; + NIR_PASS(progress, ordered_shaders[i], nir_lower_global_vars_to_local); + if (progress) { ac_nir_lower_indirect_derefs(ordered_shaders[i], pipeline->device->physical_device->rad_info.gfx_level); /* remove dead writes, which can remove input loads */ - nir_lower_vars_to_ssa(ordered_shaders[i]); - nir_opt_dce(ordered_shaders[i]); + NIR_PASS(_, ordered_shaders[i], nir_lower_vars_to_ssa); + NIR_PASS(_, ordered_shaders[i], nir_opt_dce); } - if (nir_lower_global_vars_to_local(ordered_shaders[i - 1])) { + progress = false; + NIR_PASS(progress, ordered_shaders[i - 1], nir_lower_global_vars_to_local); + if (progress) { ac_nir_lower_indirect_derefs(ordered_shaders[i - 1], pipeline->device->physical_device->rad_info.gfx_level); } @@ -4518,7 +4524,7 @@ radv_create_shaders(struct radv_pipeline *pipeline, struct radv_pipeline_layout assert(pipeline->graphics.last_vgt_api_stage == MESA_SHADER_VERTEX || pipeline->graphics.last_vgt_api_stage == MESA_SHADER_GEOMETRY); nir_shader *last_vgt_shader = stages[pipeline->graphics.last_vgt_api_stage].nir; - NIR_PASS_V(last_vgt_shader, radv_force_primitive_shading_rate, device); + NIR_PASS(_, last_vgt_shader, radv_force_primitive_shading_rate, device); } bool optimize_conservatively = pipeline_key->optimisations_disabled; @@ -4547,12 +4553,12 @@ radv_create_shaders(struct radv_pipeline *pipeline, struct radv_pipeline_layout } if (stages[MESA_SHADER_VERTEX].nir) { - NIR_PASS_V(stages[MESA_SHADER_VERTEX].nir, radv_lower_vs_input, pipeline_key); + NIR_PASS(_, stages[MESA_SHADER_VERTEX].nir, radv_lower_vs_input, pipeline_key); } if (stages[MESA_SHADER_FRAGMENT].nir && !radv_use_llvm_for_stage(device, MESA_SHADER_FRAGMENT)) { /* TODO: Convert the LLVM backend. */ - NIR_PASS_V(stages[MESA_SHADER_FRAGMENT].nir, radv_lower_fs_output, pipeline_key); + NIR_PASS(_, stages[MESA_SHADER_FRAGMENT].nir, radv_lower_fs_output, pipeline_key); } radv_fill_shader_info(pipeline, pipeline_layout, pipeline_key, stages); @@ -4594,8 +4600,8 @@ radv_create_shaders(struct radv_pipeline *pipeline, struct radv_pipeline_layout radv_declare_pipeline_args(device, stages, pipeline_key); if (stages[MESA_SHADER_FRAGMENT].nir) { - NIR_PASS_V(stages[MESA_SHADER_FRAGMENT].nir, radv_lower_fs_intrinsics, - &stages[MESA_SHADER_FRAGMENT], pipeline_key); + NIR_PASS(_, stages[MESA_SHADER_FRAGMENT].nir, radv_lower_fs_intrinsics, + &stages[MESA_SHADER_FRAGMENT], pipeline_key); } for (int i = 0; i < MESA_VULKAN_SHADER_STAGES; ++i) { @@ -4612,9 +4618,9 @@ radv_create_shaders(struct radv_pipeline *pipeline, struct radv_pipeline_layout .callback = &non_uniform_access_callback, .callback_data = NULL, }; - NIR_PASS_V(stages[i].nir, nir_lower_non_uniform_access, &options); + NIR_PASS(_, stages[i].nir, nir_lower_non_uniform_access, &options); } - NIR_PASS_V(stages[i].nir, nir_lower_memory_model); + NIR_PASS(_, stages[i].nir, nir_lower_memory_model); nir_load_store_vectorize_options vectorize_opts = { .modes = nir_var_mem_ssbo | nir_var_mem_ubo | nir_var_mem_push_const | @@ -4632,9 +4638,12 @@ radv_create_shaders(struct radv_pipeline *pipeline, struct radv_pipeline_layout nir_var_mem_ubo | nir_var_mem_ssbo | nir_var_mem_push_const; } - if (nir_opt_load_store_vectorize(stages[i].nir, &vectorize_opts)) { - NIR_PASS_V(stages[i].nir, nir_copy_prop); - nir_opt_shrink_stores(stages[i].nir, !device->instance->disable_shrink_image_store); + bool progress = false; + NIR_PASS(progress, stages[i].nir, nir_opt_load_store_vectorize, &vectorize_opts); + if (progress) { + NIR_PASS(_, stages[i].nir, nir_copy_prop); + NIR_PASS(_, stages[i].nir, nir_opt_shrink_stores, + !device->instance->disable_shrink_image_store); /* Gather info again, to update whether 8/16-bit are used. */ nir_shader_gather_info(stages[i].nir, nir_shader_get_entrypoint(stages[i].nir)); @@ -4649,31 +4658,32 @@ radv_create_shaders(struct radv_pipeline *pipeline, struct radv_pipeline_layout else if (i == MESA_SHADER_TESS_EVAL && stages[MESA_SHADER_GEOMETRY].nir) info = &stages[MESA_SHADER_GEOMETRY].info; } - NIR_PASS_V(stages[i].nir, radv_nir_lower_ycbcr_textures, pipeline_layout); + NIR_PASS(_, stages[i].nir, radv_nir_lower_ycbcr_textures, pipeline_layout); NIR_PASS_V(stages[i].nir, radv_nir_apply_pipeline_layout, device, pipeline_layout, info, &stages[i].args); - nir_opt_shrink_vectors(stages[i].nir); + NIR_PASS(_, stages[i].nir, nir_opt_shrink_vectors); - nir_lower_alu_to_scalar(stages[i].nir, NULL, NULL); + NIR_PASS(_, stages[i].nir, nir_lower_alu_to_scalar, NULL, NULL); /* lower ALU operations */ - nir_lower_int64(stages[i].nir); + NIR_PASS(_, stages[i].nir, nir_lower_int64); - nir_opt_idiv_const(stages[i].nir, 8); + NIR_PASS(_, stages[i].nir, nir_opt_idiv_const, 8); - nir_lower_idiv(stages[i].nir, - &(nir_lower_idiv_options){ - .imprecise_32bit_lowering = false, - .allow_fp16 = device->physical_device->rad_info.gfx_level >= GFX9, - }); + NIR_PASS(_, stages[i].nir, nir_lower_idiv, + &(nir_lower_idiv_options){ + .imprecise_32bit_lowering = false, + .allow_fp16 = device->physical_device->rad_info.gfx_level >= GFX9, + }); nir_move_options sink_opts = nir_move_const_undef | nir_move_copies; if (i != MESA_SHADER_FRAGMENT || !pipeline_key->disable_sinking_load_input_fs) sink_opts |= nir_move_load_input; - nir_opt_sink(stages[i].nir, sink_opts); - nir_opt_move(stages[i].nir, nir_move_load_input | nir_move_const_undef | nir_move_copies); + NIR_PASS(_, stages[i].nir, nir_opt_sink, sink_opts); + NIR_PASS(_, stages[i].nir, nir_opt_move, + nir_move_load_input | nir_move_const_undef | nir_move_copies); /* Lower I/O intrinsics to memory instructions. */ bool io_to_mem = radv_lower_io_to_mem(device, &stages[i], pipeline_key); @@ -4682,26 +4692,26 @@ radv_create_shaders(struct radv_pipeline *pipeline, struct radv_pipeline_layout if (lowered_ngg) radv_lower_ngg(device, &stages[i], pipeline_key); - ac_nir_lower_global_access(stages[i].nir); - radv_nir_lower_abi(stages[i].nir, device->physical_device->rad_info.gfx_level, - &stages[i].info, &stages[i].args, pipeline_key, - radv_use_llvm_for_stage(device, i)); + NIR_PASS(_, stages[i].nir, ac_nir_lower_global_access); + NIR_PASS_V(stages[i].nir, radv_nir_lower_abi, device->physical_device->rad_info.gfx_level, + &stages[i].info, &stages[i].args, pipeline_key, + radv_use_llvm_for_stage(device, i)); radv_optimize_nir_algebraic( stages[i].nir, io_to_mem || lowered_ngg || i == MESA_SHADER_COMPUTE || i == MESA_SHADER_TASK); if (stages[i].nir->info.bit_sizes_int & (8 | 16)) { if (device->physical_device->rad_info.gfx_level >= GFX8) { - nir_convert_to_lcssa(stages[i].nir, true, true); + NIR_PASS(_, stages[i].nir, nir_convert_to_lcssa, true, true); nir_divergence_analysis(stages[i].nir); } if (nir_lower_bit_size(stages[i].nir, lower_bit_size_callback, device)) { - NIR_PASS_V(stages[i].nir, nir_opt_constant_folding); - NIR_PASS_V(stages[i].nir, nir_opt_dce); + NIR_PASS(_, stages[i].nir, nir_opt_constant_folding); + NIR_PASS(_, stages[i].nir, nir_opt_dce); } if (device->physical_device->rad_info.gfx_level >= GFX8) - nir_opt_remove_phis(stages[i].nir); /* cleanup LCSSA phis */ + NIR_PASS(_, stages[i].nir, nir_opt_remove_phis); /* cleanup LCSSA phis */ } if (((stages[i].nir->info.bit_sizes_int | stages[i].nir->info.bit_sizes_float) & 16) && device->physical_device->rad_info.gfx_level >= GFX9) { @@ -4716,23 +4726,22 @@ radv_create_shaders(struct radv_pipeline *pipeline, struct radv_pipeline_layout NIR_PASS(copy_prop, stages[i].nir, nir_fold_16bit_image_load_store_conversions); if (copy_prop) { - NIR_PASS_V(stages[i].nir, nir_copy_prop); - NIR_PASS_V(stages[i].nir, nir_opt_dce); + NIR_PASS(_, stages[i].nir, nir_copy_prop); + NIR_PASS(_, stages[i].nir, nir_opt_dce); } - - NIR_PASS_V(stages[i].nir, nir_opt_vectorize, opt_vectorize_callback, NULL); + NIR_PASS(_, stages[i].nir, nir_opt_vectorize, opt_vectorize_callback, NULL); } /* cleanup passes */ - nir_lower_load_const_to_scalar(stages[i].nir); + NIR_PASS(_, stages[i].nir, nir_lower_load_const_to_scalar); sink_opts |= nir_move_comparisons | nir_move_load_ubo | nir_move_load_ssbo; - nir_opt_sink(stages[i].nir, sink_opts); + NIR_PASS(_, stages[i].nir, nir_opt_sink, sink_opts); nir_move_options move_opts = nir_move_const_undef | nir_move_load_ubo | nir_move_load_input | nir_move_comparisons | nir_move_copies; - nir_opt_move(stages[i].nir, move_opts); + NIR_PASS(_, stages[i].nir, nir_opt_move, move_opts); stages[i].feedback.duration += os_time_get_nano() - stage_start; } diff --git a/src/amd/vulkan/radv_pipeline_rt.c b/src/amd/vulkan/radv_pipeline_rt.c index 7b77c36c217..89ab21b820e 100644 --- a/src/amd/vulkan/radv_pipeline_rt.c +++ b/src/amd/vulkan/radv_pipeline_rt.c @@ -714,9 +714,9 @@ insert_rt_case(nir_builder *b, nir_shader *shader, const struct rt_variables *va NIR_PASS_V(shader, lower_rt_instructions, &src_vars, call_idx_base); - NIR_PASS_V(shader, nir_opt_remove_phis); - NIR_PASS_V(shader, nir_lower_returns); - NIR_PASS_V(shader, nir_opt_dce); + NIR_PASS(_, shader, nir_opt_remove_phis); + NIR_PASS(_, shader, nir_lower_returns); + NIR_PASS(_, shader, nir_opt_dce); if (b->shader->info.stage == MESA_SHADER_ANY_HIT || b->shader->info.stage == MESA_SHADER_INTERSECTION) { @@ -819,14 +819,14 @@ parse_rt_stage(struct radv_device *device, const VkPipelineShaderStageCreateInfo nir_rt_return_amd(&b_inner); } - NIR_PASS_V(shader, nir_lower_vars_to_explicit_types, - nir_var_function_temp | nir_var_shader_call_data | nir_var_ray_hit_attrib, - glsl_get_natural_size_align_bytes); + NIR_PASS(_, shader, nir_lower_vars_to_explicit_types, + nir_var_function_temp | nir_var_shader_call_data | nir_var_ray_hit_attrib, + glsl_get_natural_size_align_bytes); - NIR_PASS_V(shader, lower_rt_derefs); + NIR_PASS(_, shader, lower_rt_derefs); - NIR_PASS_V(shader, nir_lower_explicit_io, nir_var_function_temp, - nir_address_format_32bit_offset); + NIR_PASS(_, shader, nir_lower_explicit_io, nir_var_function_temp, + nir_address_format_32bit_offset); return shader; } @@ -950,7 +950,7 @@ nir_lower_intersection_shader(nir_shader *intersection, nir_shader *any_hit) struct hash_table *any_hit_var_remap = NULL; if (any_hit) { any_hit = nir_shader_clone(dead_ctx, any_hit); - NIR_PASS_V(any_hit, nir_opt_dce); + NIR_PASS(_, any_hit, nir_opt_dce); any_hit_impl = lower_any_hit_for_intersection(any_hit); any_hit_var_remap = _mesa_pointer_hash_table_create(dead_ctx); } @@ -1020,7 +1020,7 @@ nir_lower_intersection_shader(nir_shader *intersection, nir_shader *any_hit) nir_index_ssa_defs(impl); /* Eliminate the casts introduced for the commit return of the any-hit shader. */ - NIR_PASS_V(intersection, nir_opt_deref); + NIR_PASS(_, intersection, nir_opt_deref); ralloc_free(dead_ctx); } diff --git a/src/amd/vulkan/radv_shader.c b/src/amd/vulkan/radv_shader.c index ba1345297c6..b6568163ccd 100644 --- a/src/amd/vulkan/radv_shader.c +++ b/src/amd/vulkan/radv_shader.c @@ -166,10 +166,10 @@ radv_optimize_nir(struct nir_shader *shader, bool optimize_conservatively, bool NIR_PASS(progress, shader, nir_opt_copy_prop_vars); NIR_PASS(progress, shader, nir_opt_dead_write_vars); - NIR_PASS_V(shader, nir_lower_vars_to_ssa); + NIR_PASS(_, shader, nir_lower_vars_to_ssa); - NIR_PASS_V(shader, nir_lower_alu_to_scalar, NULL, NULL); - NIR_PASS_V(shader, nir_lower_phis_to_scalar, true); + NIR_PASS(_, shader, nir_lower_alu_to_scalar, NULL, NULL); + NIR_PASS(_, shader, nir_lower_phis_to_scalar, true); NIR_PASS(progress, shader, nir_copy_prop); NIR_PASS(progress, shader, nir_opt_remove_phis); @@ -207,10 +207,10 @@ radv_optimize_nir_algebraic(nir_shader *nir, bool opt_offsets) bool more_algebraic = true; while (more_algebraic) { more_algebraic = false; - NIR_PASS_V(nir, nir_copy_prop); - NIR_PASS_V(nir, nir_opt_dce); - NIR_PASS_V(nir, nir_opt_constant_folding); - NIR_PASS_V(nir, nir_opt_cse); + NIR_PASS(_, nir, nir_copy_prop); + NIR_PASS(_, nir, nir_opt_dce); + NIR_PASS(_, nir, nir_opt_constant_folding); + NIR_PASS(_, nir, nir_opt_cse); NIR_PASS(more_algebraic, nir, nir_opt_algebraic); } @@ -220,7 +220,7 @@ radv_optimize_nir_algebraic(nir_shader *nir, bool opt_offsets) .buffer_max = ~0, .shared_max = ~0, }; - NIR_PASS_V(nir, nir_opt_offsets, &offset_options); + NIR_PASS(_, nir, nir_opt_offsets, &offset_options); } /* Do late algebraic optimization to turn add(a, @@ -233,10 +233,10 @@ radv_optimize_nir_algebraic(nir_shader *nir, bool opt_offsets) 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_V(nir, nir_copy_prop); - NIR_PASS_V(nir, nir_opt_dce); - NIR_PASS_V(nir, nir_opt_cse); + NIR_PASS(_, nir, nir_opt_constant_folding); + NIR_PASS(_, nir, nir_copy_prop); + NIR_PASS(_, nir, nir_opt_dce); + NIR_PASS(_, nir, nir_opt_cse); } } @@ -712,11 +712,13 @@ radv_shader_spirv_to_nir(struct radv_device *device, const struct radv_pipeline_ * inline functions. That way they get properly initialized at the top * of the function and not at the top of its caller. */ - NIR_PASS_V(nir, nir_lower_variable_initializers, nir_var_function_temp); - NIR_PASS_V(nir, nir_lower_returns); - if (nir_inline_functions(nir)) - NIR_PASS_V(nir, nir_copy_prop); - NIR_PASS_V(nir, nir_opt_deref); + NIR_PASS(_, nir, nir_lower_variable_initializers, nir_var_function_temp); + NIR_PASS(_, nir, nir_lower_returns); + bool progress = false; + NIR_PASS(progress, nir, nir_inline_functions); + if (progress) + NIR_PASS(_, nir, nir_copy_prop); + NIR_PASS(_, nir, nir_opt_deref); /* Pick off the single entrypoint that we want */ foreach_list_typed_safe(nir_function, func, node, &nir->functions) @@ -731,43 +733,43 @@ radv_shader_spirv_to_nir(struct radv_device *device, const struct radv_pipeline_ /* Make sure we lower constant initializers on output variables so that * nir_remove_dead_variables below sees the corresponding stores */ - NIR_PASS_V(nir, nir_lower_variable_initializers, nir_var_shader_out); + NIR_PASS(_, nir, nir_lower_variable_initializers, nir_var_shader_out); /* Now that we've deleted all but the main function, we can go ahead and * lower the rest of the constant initializers. */ - NIR_PASS_V(nir, nir_lower_variable_initializers, ~0); + NIR_PASS(_, nir, nir_lower_variable_initializers, ~0); /* Split member structs. We do this before lower_io_to_temporaries so that * it doesn't lower system values to temporaries by accident. */ - NIR_PASS_V(nir, nir_split_var_copies); - NIR_PASS_V(nir, nir_split_per_member_structs); + NIR_PASS(_, nir, nir_split_var_copies); + NIR_PASS(_, nir, nir_split_per_member_structs); if (nir->info.stage == MESA_SHADER_FRAGMENT) - NIR_PASS_V(nir, nir_lower_io_to_vector, nir_var_shader_out); + NIR_PASS(_, nir, nir_lower_io_to_vector, nir_var_shader_out); if (nir->info.stage == MESA_SHADER_FRAGMENT) - NIR_PASS_V(nir, nir_lower_input_attachments, - &(nir_input_attachment_options){ - .use_fragcoord_sysval = true, - .use_layer_id_sysval = false, - }); + NIR_PASS(_, nir, nir_lower_input_attachments, + &(nir_input_attachment_options){ + .use_fragcoord_sysval = true, + .use_layer_id_sysval = false, + }); - NIR_PASS_V(nir, nir_remove_dead_variables, - nir_var_shader_in | nir_var_shader_out | nir_var_system_value | nir_var_mem_shared, - NULL); + NIR_PASS(_, nir, nir_remove_dead_variables, + nir_var_shader_in | nir_var_shader_out | nir_var_system_value | nir_var_mem_shared, + NULL); /* Variables can make nir_propagate_invariant more conservative * than it needs to be. */ - NIR_PASS_V(nir, nir_lower_global_vars_to_local); - NIR_PASS_V(nir, nir_lower_vars_to_ssa); + NIR_PASS(_, nir, nir_lower_global_vars_to_local); + NIR_PASS(_, nir, nir_lower_vars_to_ssa); - NIR_PASS_V(nir, nir_propagate_invariant, key->invariant_geom); + NIR_PASS(_, nir, nir_propagate_invariant, key->invariant_geom); - NIR_PASS_V(nir, nir_lower_clip_cull_distance_arrays); + NIR_PASS(_, nir, nir_lower_clip_cull_distance_arrays); - NIR_PASS_V(nir, nir_lower_discard_or_demote, key->ps.lower_discard_to_demote); + NIR_PASS(_, nir, nir_lower_discard_or_demote, key->ps.lower_discard_to_demote); nir_lower_doubles_options lower_doubles = nir->options->lower_doubles_options; @@ -779,10 +781,10 @@ radv_shader_spirv_to_nir(struct radv_device *device, const struct radv_pipeline_ lower_doubles |= nir_lower_dfloor; } - NIR_PASS_V(nir, nir_lower_doubles, NULL, lower_doubles); + NIR_PASS(_, nir, nir_lower_doubles, NULL, lower_doubles); } - NIR_PASS_V(nir, nir_lower_system_values); + NIR_PASS(_, nir, nir_lower_system_values); nir_lower_compute_system_values_options csv_options = { /* Mesh shaders run as NGG which can implement local_invocation_index from * the wave ID in merged_wave_info, but they don't have local_invocation_ids. @@ -793,7 +795,7 @@ radv_shader_spirv_to_nir(struct radv_device *device, const struct radv_pipeline_ (nir->info.workgroup_size[1] == 1) + (nir->info.workgroup_size[2] == 1)) == 2, }; - NIR_PASS_V(nir, nir_lower_compute_system_values, &csv_options); + NIR_PASS(_, nir, nir_lower_compute_system_values, &csv_options); /* Vulkan uses the separate-shader linking model */ nir->info.separate_shader = true; @@ -801,8 +803,8 @@ radv_shader_spirv_to_nir(struct radv_device *device, const struct radv_pipeline_ nir_shader_gather_info(nir, nir_shader_get_entrypoint(nir)); if (nir->info.ray_queries > 0) { - NIR_PASS_V(nir, nir_opt_ray_queries); - NIR_PASS_V(nir, radv_nir_lower_ray_queries, device); + NIR_PASS(_, nir, nir_opt_ray_queries); + NIR_PASS(_, nir, radv_nir_lower_ray_queries, device); } if (nir->info.stage == MESA_SHADER_GEOMETRY) { @@ -815,7 +817,7 @@ radv_shader_spirv_to_nir(struct radv_device *device, const struct radv_pipeline_ nir_lower_gs_intrinsics_overwrite_incomplete; } - nir_lower_gs_intrinsics(nir, nir_gs_flags); + NIR_PASS(_, nir, nir_lower_gs_intrinsics, nir_gs_flags); } static const nir_lower_tex_options tex_options = { @@ -827,15 +829,15 @@ radv_shader_spirv_to_nir(struct radv_device *device, const struct radv_pipeline_ .lower_invalid_implicit_lod = true, }; - nir_lower_tex(nir, &tex_options); + NIR_PASS(_, nir, nir_lower_tex, &tex_options); static const nir_lower_image_options image_options = { .lower_cube_size = true, }; - nir_lower_image(nir, &image_options); + NIR_PASS(_, nir, nir_lower_image, &image_options); - nir_lower_vars_to_ssa(nir); + NIR_PASS(_, nir, nir_lower_vars_to_ssa); if (nir->info.stage == MESA_SHADER_VERTEX || nir->info.stage == MESA_SHADER_GEOMETRY || nir->info.stage == MESA_SHADER_FRAGMENT) { @@ -844,27 +846,28 @@ radv_shader_spirv_to_nir(struct radv_device *device, const struct radv_pipeline_ NIR_PASS_V(nir, nir_lower_io_to_temporaries, nir_shader_get_entrypoint(nir), true, false); } - nir_split_var_copies(nir); + NIR_PASS(_, nir, nir_split_var_copies); - nir_lower_global_vars_to_local(nir); - nir_remove_dead_variables(nir, nir_var_function_temp, NULL); + NIR_PASS(_, nir, nir_lower_global_vars_to_local); + NIR_PASS(_, nir, nir_remove_dead_variables, nir_var_function_temp, NULL); bool gfx7minus = device->physical_device->rad_info.gfx_level <= GFX7; - nir_lower_subgroups(nir, &(struct nir_lower_subgroups_options){ - .subgroup_size = subgroup_size, - .ballot_bit_size = ballot_bit_size, - .ballot_components = 1, - .lower_to_scalar = 1, - .lower_subgroup_masks = 1, - .lower_relative_shuffle = 1, - .lower_shuffle_to_32bit = 1, - .lower_vote_eq = 1, - .lower_quad_broadcast_dynamic = 1, - .lower_quad_broadcast_dynamic_to_const = gfx7minus, - .lower_shuffle_to_swizzle_amd = 1, - }); + NIR_PASS(_, nir, nir_lower_subgroups, + &(struct nir_lower_subgroups_options){ + .subgroup_size = subgroup_size, + .ballot_bit_size = ballot_bit_size, + .ballot_components = 1, + .lower_to_scalar = 1, + .lower_subgroup_masks = 1, + .lower_relative_shuffle = 1, + .lower_shuffle_to_32bit = 1, + .lower_vote_eq = 1, + .lower_quad_broadcast_dynamic = 1, + .lower_quad_broadcast_dynamic_to_const = gfx7minus, + .lower_shuffle_to_swizzle_amd = 1, + }); - nir_lower_load_const_to_scalar(nir); - nir_opt_shrink_stores(nir, !device->instance->disable_shrink_image_store); + NIR_PASS(_, nir, nir_lower_load_const_to_scalar); + NIR_PASS(_, nir, nir_opt_shrink_stores, !device->instance->disable_shrink_image_store); if (!key->optimisations_disabled) radv_optimize_nir(nir, false, true); @@ -872,28 +875,30 @@ radv_shader_spirv_to_nir(struct radv_device *device, const struct radv_pipeline_ /* We call nir_lower_var_copies() after the first radv_optimize_nir() * to remove any copies introduced by nir_opt_find_array_copies(). */ - nir_lower_var_copies(nir); + NIR_PASS(_, nir, nir_lower_var_copies); unsigned lower_flrp = (nir->options->lower_flrp16 ? 16 : 0) | (nir->options->lower_flrp32 ? 32 : 0) | (nir->options->lower_flrp64 ? 64 : 0); if (lower_flrp != 0) { - if (nir_lower_flrp(nir, lower_flrp, false /* always_precise */)) - NIR_PASS_V(nir, nir_opt_constant_folding); + bool progress = false; + NIR_PASS(progress, nir, nir_lower_flrp, lower_flrp, false /* always precise */); + if (progress) + NIR_PASS(_, nir, nir_opt_constant_folding); } const nir_opt_access_options opt_access_options = { .is_vulkan = true, .infer_non_readable = true, }; - NIR_PASS_V(nir, nir_opt_access, &opt_access_options); + NIR_PASS(_, nir, nir_opt_access, &opt_access_options); - NIR_PASS_V(nir, nir_lower_explicit_io, nir_var_mem_push_const, nir_address_format_32bit_offset); + NIR_PASS(_, nir, nir_lower_explicit_io, nir_var_mem_push_const, nir_address_format_32bit_offset); - NIR_PASS_V(nir, nir_lower_explicit_io, nir_var_mem_ubo | nir_var_mem_ssbo, - nir_address_format_vec2_index_32bit_offset); + NIR_PASS(_, nir, nir_lower_explicit_io, nir_var_mem_ubo | nir_var_mem_ssbo, + nir_address_format_vec2_index_32bit_offset); - NIR_PASS_V(nir, lower_intrinsics, key); + NIR_PASS(_, nir, lower_intrinsics, key); /* Lower deref operations for compute shared memory. */ if (nir->info.stage == MESA_SHADER_COMPUTE || @@ -906,25 +911,25 @@ radv_shader_spirv_to_nir(struct radv_device *device, const struct radv_pipeline_ var_modes |= nir_var_mem_task_payload; if (!nir->info.shared_memory_explicit_layout) { - NIR_PASS_V(nir, nir_lower_vars_to_explicit_types, var_modes, shared_var_info); + NIR_PASS(_, nir, nir_lower_vars_to_explicit_types, var_modes, shared_var_info); } - NIR_PASS_V(nir, nir_lower_explicit_io, var_modes, nir_address_format_32bit_offset); + NIR_PASS(_, nir, nir_lower_explicit_io, var_modes, nir_address_format_32bit_offset); if (nir->info.zero_initialize_shared_memory && nir->info.shared_size > 0) { 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_zero_initialize_shared_memory, shared_size, chunk_size); + NIR_PASS(_, nir, nir_zero_initialize_shared_memory, shared_size, chunk_size); } } - nir_lower_explicit_io(nir, nir_var_mem_global | nir_var_mem_constant, - nir_address_format_64bit_global); + NIR_PASS(_, nir, nir_lower_explicit_io, nir_var_mem_global | nir_var_mem_constant, + nir_address_format_64bit_global); /* Lower large variables that are always constant with load_constant * intrinsics, which get turned into PC-relative loads from a data * section next to the shader. */ - NIR_PASS_V(nir, nir_opt_large_constants, glsl_get_natural_size_align_bytes, 16); + NIR_PASS(_, nir, nir_opt_large_constants, glsl_get_natural_size_align_bytes, 16); /* Lower primitive shading rate to match HW requirements. */ if ((nir->info.stage == MESA_SHADER_VERTEX || @@ -932,7 +937,7 @@ radv_shader_spirv_to_nir(struct radv_device *device, const struct radv_pipeline_ nir->info.stage == MESA_SHADER_MESH) && nir->info.outputs_written & BITFIELD64_BIT(VARYING_SLOT_PRIMITIVE_SHADING_RATE)) { /* Lower primitive shading rate to match HW requirements. */ - NIR_PASS_V(nir, radv_lower_primitive_shading_rate); + NIR_PASS(_, nir, radv_lower_primitive_shading_rate); } /* Indirect lowering must be called after the radv_optimize_nir() loop @@ -1025,17 +1030,17 @@ radv_lower_io(struct radv_device *device, nir_shader *nir) return; if (nir->info.stage == MESA_SHADER_FRAGMENT) { - NIR_PASS_V(nir, lower_view_index); + NIR_PASS(_, nir, lower_view_index); nir_assign_io_var_locations(nir, nir_var_shader_in, &nir->num_inputs, MESA_SHADER_FRAGMENT); } - NIR_PASS_V(nir, nir_lower_io, nir_var_shader_in | nir_var_shader_out, type_size_vec4, - nir_lower_io_lower_64bit_to_32); + NIR_PASS(_, nir, nir_lower_io, nir_var_shader_in | nir_var_shader_out, type_size_vec4, + nir_lower_io_lower_64bit_to_32); /* This pass needs actual constants */ - nir_opt_constant_folding(nir); + NIR_PASS(_, nir, nir_opt_constant_folding); - 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); } bool @@ -1047,36 +1052,36 @@ radv_lower_io_to_mem(struct radv_device *device, struct radv_pipeline_stage *sta if (nir->info.stage == MESA_SHADER_VERTEX) { if (info->vs.as_ls) { - ac_nir_lower_ls_outputs_to_mem(nir, info->vs.tcs_in_out_eq, - info->vs.tcs_temp_only_input_mask, - info->vs.num_linked_outputs); + NIR_PASS_V(nir, ac_nir_lower_ls_outputs_to_mem, info->vs.tcs_in_out_eq, + info->vs.tcs_temp_only_input_mask, info->vs.num_linked_outputs); return true; } else if (info->vs.as_es) { - ac_nir_lower_es_outputs_to_mem(nir, device->physical_device->rad_info.gfx_level, - info->vs.num_linked_outputs); + NIR_PASS_V(nir, ac_nir_lower_es_outputs_to_mem, + device->physical_device->rad_info.gfx_level, info->vs.num_linked_outputs); return true; } } else if (nir->info.stage == MESA_SHADER_TESS_CTRL) { - ac_nir_lower_hs_inputs_to_mem(nir, info->vs.tcs_in_out_eq, info->tcs.num_linked_inputs); - ac_nir_lower_hs_outputs_to_mem( - nir, device->physical_device->rad_info.gfx_level, info->tcs.tes_reads_tess_factors, - info->tcs.tes_inputs_read, info->tcs.tes_patch_inputs_read, info->tcs.num_linked_inputs, - info->tcs.num_linked_outputs, info->tcs.num_linked_patch_outputs, true); + NIR_PASS_V(nir, ac_nir_lower_hs_inputs_to_mem, info->vs.tcs_in_out_eq, + info->tcs.num_linked_inputs); + NIR_PASS_V(nir, ac_nir_lower_hs_outputs_to_mem, device->physical_device->rad_info.gfx_level, + info->tcs.tes_reads_tess_factors, info->tcs.tes_inputs_read, + info->tcs.tes_patch_inputs_read, info->tcs.num_linked_inputs, + info->tcs.num_linked_outputs, info->tcs.num_linked_patch_outputs, true); return true; } else if (nir->info.stage == MESA_SHADER_TESS_EVAL) { - ac_nir_lower_tes_inputs_to_mem(nir, info->tes.num_linked_inputs, - info->tes.num_linked_patch_inputs); + NIR_PASS_V(nir, ac_nir_lower_tes_inputs_to_mem, info->tes.num_linked_inputs, + info->tes.num_linked_patch_inputs); if (info->tes.as_es) { - ac_nir_lower_es_outputs_to_mem(nir, device->physical_device->rad_info.gfx_level, - info->tes.num_linked_outputs); + NIR_PASS_V(nir, ac_nir_lower_es_outputs_to_mem, + device->physical_device->rad_info.gfx_level, info->tes.num_linked_outputs); } return true; } else if (nir->info.stage == MESA_SHADER_GEOMETRY) { - ac_nir_lower_gs_inputs_to_mem(nir, device->physical_device->rad_info.gfx_level, - info->gs.num_linked_inputs); + NIR_PASS_V(nir, ac_nir_lower_gs_inputs_to_mem, device->physical_device->rad_info.gfx_level, + info->gs.num_linked_inputs); return true; } else if (nir->info.stage == MESA_SHADER_TASK) { ac_nir_apply_first_task_to_task_shader(nir); @@ -1216,29 +1221,17 @@ void radv_lower_ngg(struct radv_device *device, struct radv_pipeline_stage *ngg_ export_prim_id = info->tes.outinfo.export_prim_id; } - ac_nir_lower_ngg_nogs( - nir, - max_vtx_in, - num_vertices_per_prim, - info->workgroup_size, - info->wave_size, - info->has_ngg_culling, - info->has_ngg_early_prim_export, - info->is_ngg_passthrough, - export_prim_id, - pl_key->vs.provoking_vtx_last, - false, - pl_key->vs.instance_rate_inputs); + NIR_PASS_V(nir, ac_nir_lower_ngg_nogs, max_vtx_in, num_vertices_per_prim, + info->workgroup_size, info->wave_size, info->has_ngg_culling, + info->has_ngg_early_prim_export, info->is_ngg_passthrough, export_prim_id, + pl_key->vs.provoking_vtx_last, false, pl_key->vs.instance_rate_inputs); } else if (nir->info.stage == MESA_SHADER_GEOMETRY) { assert(info->is_ngg); - ac_nir_lower_ngg_gs( - nir, info->wave_size, info->workgroup_size, - info->ngg_info.esgs_ring_size, - info->gs.gsvs_vertex_size, - info->ngg_info.ngg_emit_size * 4u, - pl_key->vs.provoking_vtx_last); + NIR_PASS_V(nir, ac_nir_lower_ngg_gs, info->wave_size, info->workgroup_size, + info->ngg_info.esgs_ring_size, info->gs.gsvs_vertex_size, + info->ngg_info.ngg_emit_size * 4u, pl_key->vs.provoking_vtx_last); } else if (nir->info.stage == MESA_SHADER_MESH) { - ac_nir_lower_ngg_ms(nir, info->wave_size); + NIR_PASS_V(nir, ac_nir_lower_ngg_ms, info->wave_size); } else { unreachable("invalid SW stage passed to radv_lower_ngg"); }