OpenHarmony/third_party_mesa3d
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

radv: add missing NIR_PASS() and switch from NIR_PASS_V()

Browse Source 
Unlike NIR_PASS_V(), NIR_PASS() can skip printing the shader when
NIR_DEBUG=print.

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Gitlab: https://gitlab.freedesktop.org/mesa/mesa/-/issues/5244
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/12448>
This commit is contained in:
Rhys Perry
2021-08-18 13:50:49 +01:00
committed by Marge Bot
parent d98b7817fc
commit 75c80be484
4 changed files with 200 additions and 198 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
src/amd/common/ac_nir.c
View File

@@ -74,6 +74,6 @@ ac_nir_lower_indirect_derefs(nir_shader *shader,
*/ */
indirect_mask |= nir_var_function_temp; 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; return progress;
} }

145
src/amd/vulkan/radv_pipeline.c
View File

@@ -2840,18 +2840,20 @@ radv_link_shaders(struct radv_pipeline *pipeline,
if (ordered_shaders[i]->info.stage != last) if (ordered_shaders[i]->info.stage != last)
mask = mask | nir_var_shader_out; 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 */ /* Optimize the new vector code and then remove dead vars */
nir_copy_prop(ordered_shaders[i]); NIR_PASS(_, ordered_shaders[i], nir_copy_prop);
nir_opt_shrink_vectors(ordered_shaders[i]); NIR_PASS(_, ordered_shaders[i], nir_opt_shrink_vectors);
if (ordered_shaders[i]->info.stage != last) { if (ordered_shaders[i]->info.stage != last) {
/* Optimize swizzled movs of load_const for /* Optimize swizzled movs of load_const for
* nir_link_opt_varyings's constant propagation * 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 */ /* 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 /* 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. * not have worked because the outputs were vector.
*/ */
if (ordered_shaders[i]->info.stage == MESA_SHADER_TESS_CTRL) 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_PASS(_, ordered_shaders[i], nir_opt_dce);
nir_remove_dead_variables( NIR_PASS(_, ordered_shaders[i], nir_remove_dead_variables,
ordered_shaders[i], nir_var_function_temp | nir_var_shader_in | nir_var_shader_out, nir_var_function_temp | nir_var_shader_in | nir_var_shader_out, NULL);
NULL);
} }
} }
} }
@@ -2903,9 +2904,10 @@ radv_link_shaders(struct radv_pipeline *pipeline,
} }
} }
if (fixup_derefs) { if (fixup_derefs) {
nir_fixup_deref_modes(ordered_shaders[i]); NIR_PASS_V(ordered_shaders[i], nir_fixup_deref_modes);
nir_remove_dead_variables(ordered_shaders[i], nir_var_shader_temp, NULL); NIR_PASS(_, ordered_shaders[i], nir_remove_dead_variables, nir_var_shader_temp,
nir_opt_dce(ordered_shaders[i]); NULL);
NIR_PASS(_, ordered_shaders[i], nir_opt_dce);
} }
continue; continue;
} }
@@ -2935,9 +2937,9 @@ radv_link_shaders(struct radv_pipeline *pipeline,
psiz_var->data.mode = nir_var_shader_temp; psiz_var->data.mode = nir_var_shader_temp;
info->outputs_written &= ~VARYING_BIT_PSIZ; info->outputs_written &= ~VARYING_BIT_PSIZ;
nir_fixup_deref_modes(ordered_shaders[i]); NIR_PASS_V(ordered_shaders[i], nir_fixup_deref_modes);
nir_remove_dead_variables(ordered_shaders[i], nir_var_shader_temp, NULL); NIR_PASS(_, ordered_shaders[i], nir_remove_dead_variables, nir_var_shader_temp, NULL);
nir_opt_dce(ordered_shaders[i]); 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 && if (stages[MESA_SHADER_FRAGMENT].nir &&
(stages[MESA_SHADER_FRAGMENT].nir->info.inputs_read & VARYING_BIT_VIEWPORT) && (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)) { !(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. */ /* 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 & !(stages[pipeline->graphics.last_vgt_api_stage].nir->info.outputs_written &
VARYING_BIT_LAYER)) { VARYING_BIT_LAYER)) {
nir_shader *last_vgt_shader = stages[pipeline->graphics.last_vgt_api_stage].nir; 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) { for (int i = 1; !optimize_conservatively && (i < shader_count); ++i) {
if (nir_link_opt_varyings(ordered_shaders[i], ordered_shaders[i - 1])) { if (nir_link_opt_varyings(ordered_shaders[i], ordered_shaders[i - 1])) {
nir_opt_constant_folding(ordered_shaders[i - 1]); NIR_PASS(_, ordered_shaders[i - 1], nir_opt_constant_folding);
nir_opt_algebraic(ordered_shaders[i - 1]); NIR_PASS(_, ordered_shaders[i - 1], nir_opt_algebraic);
nir_opt_dce(ordered_shaders[i - 1]); NIR_PASS(_, ordered_shaders[i - 1], nir_opt_dce);
} }
nir_remove_dead_variables(ordered_shaders[i], nir_var_shader_out, NULL); NIR_PASS(_, ordered_shaders[i], nir_remove_dead_variables, nir_var_shader_out, NULL);
nir_remove_dead_variables(ordered_shaders[i - 1], nir_var_shader_in, 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]); 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_MESH ||
(ordered_shaders[i]->info.stage == MESA_SHADER_VERTEX && has_geom_tess) || (ordered_shaders[i]->info.stage == MESA_SHADER_VERTEX && has_geom_tess) ||
(ordered_shaders[i]->info.stage == MESA_SHADER_TESS_EVAL && merged_gs)) { (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) if (ordered_shaders[i]->info.stage == MESA_SHADER_TESS_CTRL)
nir_vectorize_tess_levels(ordered_shaders[i]); NIR_PASS(_, ordered_shaders[i], nir_vectorize_tess_levels);
nir_opt_combine_stores(ordered_shaders[i], nir_var_shader_out); NIR_PASS(_, ordered_shaders[i], nir_opt_combine_stores, nir_var_shader_out);
} }
if (ordered_shaders[i - 1]->info.stage == MESA_SHADER_GEOMETRY || 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_CTRL ||
ordered_shaders[i - 1]->info.stage == MESA_SHADER_TESS_EVAL) { 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 (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], ac_nir_lower_indirect_derefs(ordered_shaders[i],
pipeline->device->physical_device->rad_info.gfx_level); pipeline->device->physical_device->rad_info.gfx_level);
/* remove dead writes, which can remove input loads */ /* remove dead writes, which can remove input loads */
nir_lower_vars_to_ssa(ordered_shaders[i]); NIR_PASS(_, ordered_shaders[i], nir_lower_vars_to_ssa);
nir_opt_dce(ordered_shaders[i]); 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], ac_nir_lower_indirect_derefs(ordered_shaders[i - 1],
pipeline->device->physical_device->rad_info.gfx_level); 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 || assert(pipeline->graphics.last_vgt_api_stage == MESA_SHADER_VERTEX ||
pipeline->graphics.last_vgt_api_stage == MESA_SHADER_GEOMETRY); pipeline->graphics.last_vgt_api_stage == MESA_SHADER_GEOMETRY);
nir_shader *last_vgt_shader = stages[pipeline->graphics.last_vgt_api_stage].nir; 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; 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) { 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)) { if (stages[MESA_SHADER_FRAGMENT].nir && !radv_use_llvm_for_stage(device, MESA_SHADER_FRAGMENT)) {
/* TODO: Convert the LLVM backend. */ /* 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); 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); radv_declare_pipeline_args(device, stages, pipeline_key);
if (stages[MESA_SHADER_FRAGMENT].nir) { if (stages[MESA_SHADER_FRAGMENT].nir) {
NIR_PASS_V(stages[MESA_SHADER_FRAGMENT].nir, radv_lower_fs_intrinsics, NIR_PASS(_, stages[MESA_SHADER_FRAGMENT].nir, radv_lower_fs_intrinsics,
&stages[MESA_SHADER_FRAGMENT], pipeline_key); &stages[MESA_SHADER_FRAGMENT], pipeline_key);
} }
for (int i = 0; i < MESA_VULKAN_SHADER_STAGES; ++i) { 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 = &non_uniform_access_callback,
.callback_data = NULL, .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 = { nir_load_store_vectorize_options vectorize_opts = {
.modes = nir_var_mem_ssbo | nir_var_mem_ubo | nir_var_mem_push_const | .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; nir_var_mem_ubo | nir_var_mem_ssbo | nir_var_mem_push_const;
} }
if (nir_opt_load_store_vectorize(stages[i].nir, &vectorize_opts)) { bool progress = false;
NIR_PASS_V(stages[i].nir, nir_copy_prop); NIR_PASS(progress, stages[i].nir, nir_opt_load_store_vectorize, &vectorize_opts);
nir_opt_shrink_stores(stages[i].nir, !device->instance->disable_shrink_image_store); 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. */ /* 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)); 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) else if (i == MESA_SHADER_TESS_EVAL && stages[MESA_SHADER_GEOMETRY].nir)
info = &stages[MESA_SHADER_GEOMETRY].info; 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, NIR_PASS_V(stages[i].nir, radv_nir_apply_pipeline_layout, device, pipeline_layout, info,
&stages[i].args); &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 */ /* 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_PASS(_, stages[i].nir, nir_lower_idiv,
&(nir_lower_idiv_options){ &(nir_lower_idiv_options){
.imprecise_32bit_lowering = false, .imprecise_32bit_lowering = false,
.allow_fp16 = device->physical_device->rad_info.gfx_level >= GFX9, .allow_fp16 = device->physical_device->rad_info.gfx_level >= GFX9,
}); });
nir_move_options sink_opts = nir_move_const_undef | nir_move_copies; nir_move_options sink_opts = nir_move_const_undef | nir_move_copies;
if (i != MESA_SHADER_FRAGMENT || !pipeline_key->disable_sinking_load_input_fs) if (i != MESA_SHADER_FRAGMENT || !pipeline_key->disable_sinking_load_input_fs)
sink_opts |= nir_move_load_input; sink_opts |= nir_move_load_input;
nir_opt_sink(stages[i].nir, sink_opts); NIR_PASS(_, stages[i].nir, nir_opt_sink, 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_move,
nir_move_load_input | nir_move_const_undef | nir_move_copies);
/* Lower I/O intrinsics to memory instructions. */ /* Lower I/O intrinsics to memory instructions. */
bool io_to_mem = radv_lower_io_to_mem(device, &stages[i], pipeline_key); 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) if (lowered_ngg)
radv_lower_ngg(device, &stages[i], pipeline_key); radv_lower_ngg(device, &stages[i], pipeline_key);
ac_nir_lower_global_access(stages[i].nir); NIR_PASS(_, stages[i].nir, ac_nir_lower_global_access);
radv_nir_lower_abi(stages[i].nir, device->physical_device->rad_info.gfx_level, 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, &stages[i].info, &stages[i].args, pipeline_key,
radv_use_llvm_for_stage(device, i)); radv_use_llvm_for_stage(device, i));
radv_optimize_nir_algebraic( radv_optimize_nir_algebraic(
stages[i].nir, io_to_mem || lowered_ngg || i == MESA_SHADER_COMPUTE || i == MESA_SHADER_TASK); 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 (stages[i].nir->info.bit_sizes_int & (8 | 16)) {
if (device->physical_device->rad_info.gfx_level >= GFX8) { 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); nir_divergence_analysis(stages[i].nir);
} }
if (nir_lower_bit_size(stages[i].nir, lower_bit_size_callback, device)) { 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(_, stages[i].nir, nir_opt_constant_folding);
NIR_PASS_V(stages[i].nir, nir_opt_dce); NIR_PASS(_, stages[i].nir, nir_opt_dce);
} }
if (device->physical_device->rad_info.gfx_level >= GFX8) 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) && if (((stages[i].nir->info.bit_sizes_int | stages[i].nir->info.bit_sizes_float) & 16) &&
device->physical_device->rad_info.gfx_level >= GFX9) { 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); NIR_PASS(copy_prop, stages[i].nir, nir_fold_16bit_image_load_store_conversions);
if (copy_prop) { if (copy_prop) {
NIR_PASS_V(stages[i].nir, nir_copy_prop); NIR_PASS(_, stages[i].nir, nir_copy_prop);
NIR_PASS_V(stages[i].nir, nir_opt_dce); NIR_PASS(_, stages[i].nir, nir_opt_dce);
} }
NIR_PASS(_, stages[i].nir, nir_opt_vectorize, opt_vectorize_callback, NULL);
NIR_PASS_V(stages[i].nir, nir_opt_vectorize, opt_vectorize_callback, NULL);
} }
/* cleanup passes */ /* 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; 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_options move_opts = nir_move_const_undef | nir_move_load_ubo |
nir_move_load_input | nir_move_comparisons | nir_move_copies; 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; stages[i].feedback.duration += os_time_get_nano() - stage_start;
} }

22
src/amd/vulkan/radv_pipeline_rt.c
View File

@@ -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, lower_rt_instructions, &src_vars, call_idx_base);
NIR_PASS_V(shader, nir_opt_remove_phis); NIR_PASS(_, shader, nir_opt_remove_phis);
NIR_PASS_V(shader, nir_lower_returns); NIR_PASS(_, shader, nir_lower_returns);
NIR_PASS_V(shader, nir_opt_dce); NIR_PASS(_, shader, nir_opt_dce);
if (b->shader->info.stage == MESA_SHADER_ANY_HIT || if (b->shader->info.stage == MESA_SHADER_ANY_HIT ||
b->shader->info.stage == MESA_SHADER_INTERSECTION) { 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_rt_return_amd(&b_inner);
} }
NIR_PASS_V(shader, nir_lower_vars_to_explicit_types, NIR_PASS(_, shader, nir_lower_vars_to_explicit_types,
nir_var_function_temp | nir_var_shader_call_data | nir_var_ray_hit_attrib, nir_var_function_temp | nir_var_shader_call_data | nir_var_ray_hit_attrib,
glsl_get_natural_size_align_bytes); 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_PASS(_, shader, nir_lower_explicit_io, nir_var_function_temp,
nir_address_format_32bit_offset); nir_address_format_32bit_offset);
return shader; 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; struct hash_table *any_hit_var_remap = NULL;
if (any_hit) { if (any_hit) {
any_hit = nir_shader_clone(dead_ctx, 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_impl = lower_any_hit_for_intersection(any_hit);
any_hit_var_remap = _mesa_pointer_hash_table_create(dead_ctx); 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); nir_index_ssa_defs(impl);
/* Eliminate the casts introduced for the commit return of the any-hit shader. */ /* 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); ralloc_free(dead_ctx);
} }

229
src/amd/vulkan/radv_shader.c
View File

@@ -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_copy_prop_vars);
NIR_PASS(progress, shader, nir_opt_dead_write_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(_, shader, nir_lower_alu_to_scalar, NULL, NULL);
NIR_PASS_V(shader, nir_lower_phis_to_scalar, true); NIR_PASS(_, shader, nir_lower_phis_to_scalar, true);
NIR_PASS(progress, shader, nir_copy_prop); NIR_PASS(progress, shader, nir_copy_prop);
NIR_PASS(progress, shader, nir_opt_remove_phis); 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; bool more_algebraic = true;
while (more_algebraic) { while (more_algebraic) {
more_algebraic = false; more_algebraic = false;
NIR_PASS_V(nir, nir_copy_prop); NIR_PASS(_, nir, nir_copy_prop);
NIR_PASS_V(nir, nir_opt_dce); NIR_PASS(_, nir, nir_opt_dce);
NIR_PASS_V(nir, nir_opt_constant_folding); NIR_PASS(_, nir, nir_opt_constant_folding);
NIR_PASS_V(nir, nir_opt_cse); NIR_PASS(_, nir, nir_opt_cse);
NIR_PASS(more_algebraic, nir, nir_opt_algebraic); 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, .buffer_max = ~0,
.shared_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, /* 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) { while (more_late_algebraic) {
more_late_algebraic = false; more_late_algebraic = false;
NIR_PASS(more_late_algebraic, nir, nir_opt_algebraic_late); 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);
NIR_PASS_V(nir, nir_copy_prop); NIR_PASS(_, nir, nir_copy_prop);
NIR_PASS_V(nir, nir_opt_dce); NIR_PASS(_, nir, nir_opt_dce);
NIR_PASS_V(nir, nir_opt_cse); 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 * inline functions. That way they get properly initialized at the top
* of the function and not at the top of its caller. * 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(_, nir, nir_lower_variable_initializers, nir_var_function_temp);
NIR_PASS_V(nir, nir_lower_returns); NIR_PASS(_, nir, nir_lower_returns);
if (nir_inline_functions(nir)) bool progress = false;
NIR_PASS_V(nir, nir_copy_prop); NIR_PASS(progress, nir, nir_inline_functions);
NIR_PASS_V(nir, nir_opt_deref); if (progress)
NIR_PASS(_, nir, nir_copy_prop);
NIR_PASS(_, nir, nir_opt_deref);
/* Pick off the single entrypoint that we want */ /* Pick off the single entrypoint that we want */
foreach_list_typed_safe(nir_function, func, node, &nir->functions) 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 /* Make sure we lower constant initializers on output variables so that
* nir_remove_dead_variables below sees the corresponding stores * 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 /* Now that we've deleted all but the main function, we can go ahead and
* lower the rest of the constant initializers. * 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 /* Split member structs. We do this before lower_io_to_temporaries so that
* it doesn't lower system values to temporaries by accident. * it doesn't lower system values to temporaries by accident.
*/ */
NIR_PASS_V(nir, nir_split_var_copies); NIR_PASS(_, nir, nir_split_var_copies);
NIR_PASS_V(nir, nir_split_per_member_structs); NIR_PASS(_, nir, nir_split_per_member_structs);
if (nir->info.stage == MESA_SHADER_FRAGMENT) 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) if (nir->info.stage == MESA_SHADER_FRAGMENT)
NIR_PASS_V(nir, nir_lower_input_attachments, NIR_PASS(_, nir, nir_lower_input_attachments,
&(nir_input_attachment_options){ &(nir_input_attachment_options){
.use_fragcoord_sysval = true, .use_fragcoord_sysval = true,
.use_layer_id_sysval = false, .use_layer_id_sysval = false,
}); });
NIR_PASS_V(nir, nir_remove_dead_variables, NIR_PASS(_, nir, nir_remove_dead_variables,
nir_var_shader_in | nir_var_shader_out | nir_var_system_value | nir_var_mem_shared, nir_var_shader_in | nir_var_shader_out | nir_var_system_value | nir_var_mem_shared,
NULL); NULL);
/* Variables can make nir_propagate_invariant more conservative /* Variables can make nir_propagate_invariant more conservative
* than it needs to be. * than it needs to be.
*/ */
NIR_PASS_V(nir, nir_lower_global_vars_to_local); NIR_PASS(_, nir, nir_lower_global_vars_to_local);
NIR_PASS_V(nir, nir_lower_vars_to_ssa); 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; 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; 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 = { nir_lower_compute_system_values_options csv_options = {
/* Mesh shaders run as NGG which can implement local_invocation_index from /* 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. * 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[1] == 1) +
(nir->info.workgroup_size[2] == 1)) == 2, (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 */ /* Vulkan uses the separate-shader linking model */
nir->info.separate_shader = true; 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)); nir_shader_gather_info(nir, nir_shader_get_entrypoint(nir));
if (nir->info.ray_queries > 0) { if (nir->info.ray_queries > 0) {
NIR_PASS_V(nir, nir_opt_ray_queries); NIR_PASS(_, nir, nir_opt_ray_queries);
NIR_PASS_V(nir, radv_nir_lower_ray_queries, device); NIR_PASS(_, nir, radv_nir_lower_ray_queries, device);
} }
if (nir->info.stage == MESA_SHADER_GEOMETRY) { 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_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 = { 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, .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 = { static const nir_lower_image_options image_options = {
.lower_cube_size = true, .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 || if (nir->info.stage == MESA_SHADER_VERTEX || nir->info.stage == MESA_SHADER_GEOMETRY ||
nir->info.stage == MESA_SHADER_FRAGMENT) { 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_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_PASS(_, nir, nir_lower_global_vars_to_local);
nir_remove_dead_variables(nir, nir_var_function_temp, NULL); NIR_PASS(_, nir, nir_remove_dead_variables, nir_var_function_temp, NULL);
bool gfx7minus = device->physical_device->rad_info.gfx_level <= GFX7; bool gfx7minus = device->physical_device->rad_info.gfx_level <= GFX7;
nir_lower_subgroups(nir, &(struct nir_lower_subgroups_options){ NIR_PASS(_, nir, nir_lower_subgroups,
.subgroup_size = subgroup_size, &(struct nir_lower_subgroups_options){
.ballot_bit_size = ballot_bit_size, .subgroup_size = subgroup_size,
.ballot_components = 1, .ballot_bit_size = ballot_bit_size,
.lower_to_scalar = 1, .ballot_components = 1,
.lower_subgroup_masks = 1, .lower_to_scalar = 1,
.lower_relative_shuffle = 1, .lower_subgroup_masks = 1,
.lower_shuffle_to_32bit = 1, .lower_relative_shuffle = 1,
.lower_vote_eq = 1, .lower_shuffle_to_32bit = 1,
.lower_quad_broadcast_dynamic = 1, .lower_vote_eq = 1,
.lower_quad_broadcast_dynamic_to_const = gfx7minus, .lower_quad_broadcast_dynamic = 1,
.lower_shuffle_to_swizzle_amd = 1, .lower_quad_broadcast_dynamic_to_const = gfx7minus,
}); .lower_shuffle_to_swizzle_amd = 1,
});
nir_lower_load_const_to_scalar(nir); NIR_PASS(_, nir, nir_lower_load_const_to_scalar);
nir_opt_shrink_stores(nir, !device->instance->disable_shrink_image_store); NIR_PASS(_, nir, nir_opt_shrink_stores, !device->instance->disable_shrink_image_store);
if (!key->optimisations_disabled) if (!key->optimisations_disabled)
radv_optimize_nir(nir, false, true); 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() /* We call nir_lower_var_copies() after the first radv_optimize_nir()
* to remove any copies introduced by nir_opt_find_array_copies(). * 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) | unsigned lower_flrp = (nir->options->lower_flrp16 ? 16 : 0) |
(nir->options->lower_flrp32 ? 32 : 0) | (nir->options->lower_flrp32 ? 32 : 0) |
(nir->options->lower_flrp64 ? 64 : 0); (nir->options->lower_flrp64 ? 64 : 0);
if (lower_flrp != 0) { if (lower_flrp != 0) {
if (nir_lower_flrp(nir, lower_flrp, false /* always_precise */)) bool progress = false;
NIR_PASS_V(nir, nir_opt_constant_folding); 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 = { const nir_opt_access_options opt_access_options = {
.is_vulkan = true, .is_vulkan = true,
.infer_non_readable = 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_PASS(_, nir, nir_lower_explicit_io, nir_var_mem_ubo | nir_var_mem_ssbo,
nir_address_format_vec2_index_32bit_offset); 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. */ /* Lower deref operations for compute shared memory. */
if (nir->info.stage == MESA_SHADER_COMPUTE || 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; var_modes |= nir_var_mem_task_payload;
if (!nir->info.shared_memory_explicit_layout) { 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) { if (nir->info.zero_initialize_shared_memory && nir->info.shared_size > 0) {
const unsigned chunk_size = 16; /* max single store size */ const unsigned chunk_size = 16; /* max single store size */
const unsigned shared_size = ALIGN(nir->info.shared_size, chunk_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_PASS(_, nir, nir_lower_explicit_io, nir_var_mem_global | nir_var_mem_constant,
nir_address_format_64bit_global); nir_address_format_64bit_global);
/* Lower large variables that are always constant with load_constant /* Lower large variables that are always constant with load_constant
* intrinsics, which get turned into PC-relative loads from a data * intrinsics, which get turned into PC-relative loads from a data
* section next to the shader. * 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. */ /* Lower primitive shading rate to match HW requirements. */
if ((nir->info.stage == MESA_SHADER_VERTEX || 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.stage == MESA_SHADER_MESH) &&
nir->info.outputs_written & BITFIELD64_BIT(VARYING_SLOT_PRIMITIVE_SHADING_RATE)) { nir->info.outputs_written & BITFIELD64_BIT(VARYING_SLOT_PRIMITIVE_SHADING_RATE)) {
/* Lower primitive shading rate to match HW requirements. */ /* 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 /* 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; return;
if (nir->info.stage == MESA_SHADER_FRAGMENT) { 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_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_PASS(_, nir, nir_lower_io, nir_var_shader_in | nir_var_shader_out, type_size_vec4,
nir_lower_io_lower_64bit_to_32); nir_lower_io_lower_64bit_to_32);
/* This pass needs actual constants */ /* 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 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 (nir->info.stage == MESA_SHADER_VERTEX) {
if (info->vs.as_ls) { if (info->vs.as_ls) {
ac_nir_lower_ls_outputs_to_mem(nir, info->vs.tcs_in_out_eq, 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.tcs_temp_only_input_mask, info->vs.num_linked_outputs);
info->vs.num_linked_outputs);
return true; return true;
} else if (info->vs.as_es) { } else if (info->vs.as_es) {
ac_nir_lower_es_outputs_to_mem(nir, device->physical_device->rad_info.gfx_level, NIR_PASS_V(nir, ac_nir_lower_es_outputs_to_mem,
info->vs.num_linked_outputs); device->physical_device->rad_info.gfx_level, info->vs.num_linked_outputs);
return true; return true;
} }
} else if (nir->info.stage == MESA_SHADER_TESS_CTRL) { } 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); NIR_PASS_V(nir, ac_nir_lower_hs_inputs_to_mem, info->vs.tcs_in_out_eq,
ac_nir_lower_hs_outputs_to_mem( info->tcs.num_linked_inputs);
nir, device->physical_device->rad_info.gfx_level, info->tcs.tes_reads_tess_factors, NIR_PASS_V(nir, ac_nir_lower_hs_outputs_to_mem, device->physical_device->rad_info.gfx_level,
info->tcs.tes_inputs_read, info->tcs.tes_patch_inputs_read, info->tcs.num_linked_inputs, info->tcs.tes_reads_tess_factors, info->tcs.tes_inputs_read,
info->tcs.num_linked_outputs, info->tcs.num_linked_patch_outputs, true); 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; return true;
} else if (nir->info.stage == MESA_SHADER_TESS_EVAL) { } else if (nir->info.stage == MESA_SHADER_TESS_EVAL) {
ac_nir_lower_tes_inputs_to_mem(nir, info->tes.num_linked_inputs, NIR_PASS_V(nir, ac_nir_lower_tes_inputs_to_mem, info->tes.num_linked_inputs,
info->tes.num_linked_patch_inputs); info->tes.num_linked_patch_inputs);
if (info->tes.as_es) { if (info->tes.as_es) {
ac_nir_lower_es_outputs_to_mem(nir, device->physical_device->rad_info.gfx_level, NIR_PASS_V(nir, ac_nir_lower_es_outputs_to_mem,
info->tes.num_linked_outputs); device->physical_device->rad_info.gfx_level, info->tes.num_linked_outputs);
} }
return true; return true;
} else if (nir->info.stage == MESA_SHADER_GEOMETRY) { } else if (nir->info.stage == MESA_SHADER_GEOMETRY) {
ac_nir_lower_gs_inputs_to_mem(nir, device->physical_device->rad_info.gfx_level, NIR_PASS_V(nir, ac_nir_lower_gs_inputs_to_mem, device->physical_device->rad_info.gfx_level,
info->gs.num_linked_inputs); info->gs.num_linked_inputs);
return true; return true;
} else if (nir->info.stage == MESA_SHADER_TASK) { } else if (nir->info.stage == MESA_SHADER_TASK) {
ac_nir_apply_first_task_to_task_shader(nir); 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; export_prim_id = info->tes.outinfo.export_prim_id;
} }
ac_nir_lower_ngg_nogs( NIR_PASS_V(nir, ac_nir_lower_ngg_nogs, max_vtx_in, num_vertices_per_prim,
nir, info->workgroup_size, info->wave_size, info->has_ngg_culling,
max_vtx_in, info->has_ngg_early_prim_export, info->is_ngg_passthrough, export_prim_id,
num_vertices_per_prim, pl_key->vs.provoking_vtx_last, false, pl_key->vs.instance_rate_inputs);
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) { } else if (nir->info.stage == MESA_SHADER_GEOMETRY) {
assert(info->is_ngg); assert(info->is_ngg);
ac_nir_lower_ngg_gs( NIR_PASS_V(nir, ac_nir_lower_ngg_gs, info->wave_size, info->workgroup_size,
nir, info->wave_size, info->workgroup_size, info->ngg_info.esgs_ring_size, info->gs.gsvs_vertex_size,
info->ngg_info.esgs_ring_size, info->ngg_info.ngg_emit_size * 4u, pl_key->vs.provoking_vtx_last);
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) { } 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 { } else {
unreachable("invalid SW stage passed to radv_lower_ngg"); unreachable("invalid SW stage passed to radv_lower_ngg");
} }