Module: Mesa Branch: main Commit: 6e1a73a53d1640c1732b4566683b8455143b0aa9 URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=6e1a73a53d1640c1732b4566683b8455143b0aa9
Author: Samuel Pitoiset <samuel.pitoi...@gmail.com> Date: Tue Jan 9 15:35:43 2024 +0100 radv: add optimisations_disabled to radv_shader_stage_key At some point, we will probably have a VK_SHADER_CREATE_xxx flag matching the pipeline one. So, I think it's more like a per-shader field. It can also be useful to disable optimisations per stage when debugging. Signed-off-by: Samuel Pitoiset <samuel.pitoi...@gmail.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27007> --- src/amd/vulkan/radv_aco_shader_info.h | 4 ++-- src/amd/vulkan/radv_pipeline.c | 20 +++++++++++--------- src/amd/vulkan/radv_pipeline_compute.c | 2 +- src/amd/vulkan/radv_pipeline_graphics.c | 33 +++++++++++++++++---------------- src/amd/vulkan/radv_pipeline_rt.c | 4 ++-- src/amd/vulkan/radv_shader.c | 27 ++++++++++++++++----------- src/amd/vulkan/radv_shader.h | 3 ++- 7 files changed, 51 insertions(+), 42 deletions(-) diff --git a/src/amd/vulkan/radv_aco_shader_info.h b/src/amd/vulkan/radv_aco_shader_info.h index 05cc681b03b..2b6be698d42 100644 --- a/src/amd/vulkan/radv_aco_shader_info.h +++ b/src/amd/vulkan/radv_aco_shader_info.h @@ -126,7 +126,7 @@ radv_aco_convert_ps_epilog_key(struct aco_ps_epilog_info *aco_info, const struct static inline void radv_aco_convert_opts(struct aco_compiler_options *aco_info, const struct radv_nir_compiler_options *radv, - const struct radv_shader_args *radv_args) + const struct radv_shader_args *radv_args, const struct radv_shader_stage_key *stage_key) { ASSIGN_FIELD(dump_shader); ASSIGN_FIELD(dump_preoptir); @@ -139,7 +139,7 @@ radv_aco_convert_opts(struct aco_compiler_options *aco_info, const struct radv_n ASSIGN_FIELD(debug.private_data); aco_info->is_opengl = false; aco_info->load_grid_size_from_user_sgpr = radv_args->load_grid_size_from_user_sgpr; - aco_info->optimisations_disabled = radv->key.optimisations_disabled; + aco_info->optimisations_disabled = stage_key->optimisations_disabled; aco_info->gfx_level = radv->info->gfx_level; aco_info->family = radv->info->family; aco_info->address32_hi = radv->info->address32_hi; diff --git a/src/amd/vulkan/radv_pipeline.c b/src/amd/vulkan/radv_pipeline.c index 3e5cde9c794..a515f063758 100644 --- a/src/amd/vulkan/radv_pipeline.c +++ b/src/amd/vulkan/radv_pipeline.c @@ -151,8 +151,10 @@ radv_generate_pipeline_key(const struct radv_device *device, const VkPipelineSha memset(&key, 0, sizeof(key)); - if (flags & VK_PIPELINE_CREATE_2_DISABLE_OPTIMIZATION_BIT_KHR) - key.optimisations_disabled = 1; + for (unsigned i = 0; i < MESA_VULKAN_SHADER_STAGES; i++) { + if (flags & VK_PIPELINE_CREATE_2_DISABLE_OPTIMIZATION_BIT_KHR) + key.stage_info[i].optimisations_disabled = 1; + } for (unsigned i = 0; i < num_stages; ++i) { const VkPipelineShaderStageCreateInfo *const stage = &stages[i]; @@ -476,7 +478,7 @@ radv_postprocess_nir(struct radv_device *device, const struct radv_pipeline_key assert(stage->info.wave_size && stage->info.workgroup_size); if (stage->stage == MESA_SHADER_FRAGMENT) { - if (!pipeline_key->optimisations_disabled) { + if (!stage->key.optimisations_disabled) { NIR_PASS(_, stage->nir, nir_opt_cse); } NIR_PASS(_, stage->nir, radv_nir_lower_fs_intrinsics, stage, pipeline_key); @@ -492,7 +494,7 @@ radv_postprocess_nir(struct radv_device *device, const struct radv_pipeline_key * thus a cheaper and likely to fail check is run first. */ if (nir_has_non_uniform_access(stage->nir, lower_non_uniform_access_types)) { - if (!pipeline_key->optimisations_disabled) { + if (!stage->key.optimisations_disabled) { NIR_PASS(_, stage->nir, nir_opt_non_uniform_access); } @@ -525,7 +527,7 @@ radv_postprocess_nir(struct radv_device *device, const struct radv_pipeline_key if (stage->key.storage_robustness2) vectorize_opts.robust_modes |= nir_var_mem_ssbo; - if (!pipeline_key->optimisations_disabled) { + if (!stage->key.optimisations_disabled) { progress = false; NIR_PASS(progress, stage->nir, nir_opt_load_store_vectorize, &vectorize_opts); if (progress) { @@ -569,7 +571,7 @@ radv_postprocess_nir(struct radv_device *device, const struct radv_pipeline_key NIR_PASS_V(stage->nir, radv_nir_apply_pipeline_layout, device, stage); - if (!pipeline_key->optimisations_disabled) { + if (!stage->key.optimisations_disabled) { NIR_PASS(_, stage->nir, nir_opt_shrink_vectors); } @@ -577,7 +579,7 @@ radv_postprocess_nir(struct radv_device *device, const struct radv_pipeline_key nir_move_options sink_opts = nir_move_const_undef | nir_move_copies; - if (!pipeline_key->optimisations_disabled) { + if (!stage->key.optimisations_disabled) { if (stage->stage != MESA_SHADER_FRAGMENT || !device->cache_key.disable_sinking_load_input_fs) sink_opts |= nir_move_load_input; @@ -718,7 +720,7 @@ radv_postprocess_nir(struct radv_device *device, const struct radv_pipeline_key }; NIR_PASS(_, stage->nir, nir_fold_16bit_tex_image, &fold_16bit_options); - if (!pipeline_key->optimisations_disabled) { + if (!stage->key.optimisations_disabled) { NIR_PASS(_, stage->nir, nir_opt_vectorize, opt_vectorize_callback, device); } } @@ -729,7 +731,7 @@ radv_postprocess_nir(struct radv_device *device, const struct radv_pipeline_key NIR_PASS(_, stage->nir, nir_copy_prop); NIR_PASS(_, stage->nir, nir_opt_dce); - if (!pipeline_key->optimisations_disabled) { + if (!stage->key.optimisations_disabled) { sink_opts |= nir_move_comparisons | nir_move_load_ubo | nir_move_load_ssbo; NIR_PASS(_, stage->nir, nir_opt_sink, sink_opts); diff --git a/src/amd/vulkan/radv_pipeline_compute.c b/src/amd/vulkan/radv_pipeline_compute.c index 7122f4055f8..1c4bf9d4946 100644 --- a/src/amd/vulkan/radv_pipeline_compute.c +++ b/src/amd/vulkan/radv_pipeline_compute.c @@ -141,7 +141,7 @@ radv_compile_cs(struct radv_device *device, struct vk_pipeline_cache *cache, str /* Compile SPIR-V shader to NIR. */ cs_stage->nir = radv_shader_spirv_to_nir(device, cs_stage, pipeline_key, is_internal); - radv_optimize_nir(cs_stage->nir, pipeline_key->optimisations_disabled); + radv_optimize_nir(cs_stage->nir, cs_stage->key.optimisations_disabled); /* Gather info again, information such as outputs_read can be out-of-date. */ nir_shader_gather_info(cs_stage->nir, nir_shader_get_entrypoint(cs_stage->nir)); diff --git a/src/amd/vulkan/radv_pipeline_graphics.c b/src/amd/vulkan/radv_pipeline_graphics.c index 874171ebbd7..f208eb8e423 100644 --- a/src/amd/vulkan/radv_pipeline_graphics.c +++ b/src/amd/vulkan/radv_pipeline_graphics.c @@ -1293,10 +1293,12 @@ merge_tess_info(struct shader_info *tes_info, struct shader_info *tcs_info) } static void -radv_link_shaders(const struct radv_device *device, nir_shader *producer, nir_shader *consumer, - const struct radv_pipeline_key *pipeline_key) +radv_link_shaders(const struct radv_device *device, struct radv_shader_stage *producer_stage, + struct radv_shader_stage *consumer_stage, const struct radv_pipeline_key *pipeline_key) { const enum amd_gfx_level gfx_level = device->physical_device->rad_info.gfx_level; + nir_shader *producer = producer_stage->nir; + nir_shader *consumer = consumer_stage->nir; bool progress; if (consumer->info.stage == MESA_SHADER_FRAGMENT) { @@ -1310,7 +1312,7 @@ radv_link_shaders(const struct radv_device *device, nir_shader *producer, nir_sh NIR_PASS(_, consumer, radv_nir_lower_view_index, producer->info.stage == MESA_SHADER_MESH); } - if (pipeline_key->optimisations_disabled) + if (producer_stage->key.optimisations_disabled || consumer_stage->key.optimisations_disabled) return; if (consumer->info.stage == MESA_SHADER_FRAGMENT && producer->info.has_transform_feedback_varyings) { @@ -1429,7 +1431,7 @@ radv_link_vs(const struct radv_device *device, struct radv_shader_stage *vs_stag next_stage->nir->info.stage == MESA_SHADER_GEOMETRY || next_stage->nir->info.stage == MESA_SHADER_FRAGMENT); - radv_link_shaders(device, vs_stage->nir, next_stage->nir, pipeline_key); + radv_link_shaders(device, vs_stage, next_stage, pipeline_key); } nir_foreach_shader_in_variable (var, vs_stage->nir) { @@ -1469,7 +1471,7 @@ radv_link_tcs(const struct radv_device *device, struct radv_shader_stage *tcs_st assert(tcs_stage->nir->info.stage == MESA_SHADER_TESS_CTRL); assert(tes_stage->nir->info.stage == MESA_SHADER_TESS_EVAL); - radv_link_shaders(device, tcs_stage->nir, tes_stage->nir, pipeline_key); + radv_link_shaders(device, tcs_stage, tes_stage, pipeline_key); /* Copy TCS info into the TES info */ merge_tess_info(&tes_stage->nir->info, &tcs_stage->nir->info); @@ -1498,7 +1500,7 @@ radv_link_tes(const struct radv_device *device, struct radv_shader_stage *tes_st assert(next_stage->nir->info.stage == MESA_SHADER_GEOMETRY || next_stage->nir->info.stage == MESA_SHADER_FRAGMENT); - radv_link_shaders(device, tes_stage->nir, next_stage->nir, pipeline_key); + radv_link_shaders(device, tes_stage, next_stage, pipeline_key); } if (next_stage && next_stage->nir->info.stage == MESA_SHADER_GEOMETRY) { @@ -1529,7 +1531,7 @@ radv_link_gs(const struct radv_device *device, struct radv_shader_stage *gs_stag if (fs_stage) { assert(fs_stage->nir->info.stage == MESA_SHADER_FRAGMENT); - radv_link_shaders(device, gs_stage->nir, fs_stage->nir, pipeline_key); + radv_link_shaders(device, gs_stage, fs_stage, pipeline_key); } nir_foreach_shader_out_variable (var, gs_stage->nir) { @@ -1545,7 +1547,7 @@ radv_link_task(const struct radv_device *device, struct radv_shader_stage *task_ assert(mesh_stage->nir->info.stage == MESA_SHADER_MESH); /* Linking task and mesh shaders shouldn't do anything for now but keep it for consistency. */ - radv_link_shaders(device, task_stage->nir, mesh_stage->nir, pipeline_key); + radv_link_shaders(device, task_stage, mesh_stage, pipeline_key); } static void @@ -1565,7 +1567,7 @@ radv_link_mesh(const struct radv_device *device, struct radv_shader_stage *mesh_ } } - radv_link_shaders(device, mesh_stage->nir, fs_stage->nir, pipeline_key); + radv_link_shaders(device, mesh_stage, fs_stage, pipeline_key); } /* ac_nir_lower_ngg ignores driver locations for mesh shaders, but set them to all zero just to @@ -2175,6 +2177,10 @@ radv_create_gs_copy_shader(struct radv_device *device, struct vk_pipeline_cache struct radv_shader_stage gs_copy_stage = { .stage = MESA_SHADER_VERTEX, .shader_sha1 = {0}, + .key = + { + .optimisations_disabled = gs_stage->key.optimisations_disabled, + }, }; radv_nir_shader_info_init(gs_copy_stage.stage, MESA_SHADER_FRAGMENT, &gs_copy_stage.info); radv_nir_shader_info_pass(device, nir, &gs_stage->layout, &gs_stage->key, pipeline_key, RADV_PIPELINE_GRAPHICS, @@ -2196,10 +2202,7 @@ radv_create_gs_copy_shader(struct radv_device *device, struct vk_pipeline_cache NIR_PASS_V(nir, radv_nir_lower_abi, device->physical_device->rad_info.gfx_level, &gs_copy_stage, pipeline_key, device->physical_device->rad_info.address32_hi); - struct radv_pipeline_key key = { - .optimisations_disabled = pipeline_key->optimisations_disabled, - }; - + struct radv_pipeline_key key = {0}; bool dump_shader = radv_can_dump_shader(device, nir, true); *gs_copy_binary = @@ -2461,8 +2464,6 @@ radv_graphics_shaders_compile(struct radv_device *device, struct vk_pipeline_cac active_nir_stages |= mesa_to_vk_shader_stage(i); } - bool optimize_conservatively = pipeline_key->optimisations_disabled; - if (!device->mesh_fast_launch_2 && stages[MESA_SHADER_MESH].nir && BITSET_TEST(stages[MESA_SHADER_MESH].nir->info.system_values_read, SYSTEM_VALUE_WORKGROUP_ID)) { nir_shader *mesh = stages[MESA_SHADER_MESH].nir; @@ -2527,7 +2528,7 @@ radv_graphics_shaders_compile(struct radv_device *device, struct vk_pipeline_cac { int64_t stage_start = os_time_get_nano(); - radv_optimize_nir(stages[i].nir, optimize_conservatively); + radv_optimize_nir(stages[i].nir, stages[i].key.optimisations_disabled); /* Gather info again, information such as outputs_read can be out-of-date. */ nir_shader_gather_info(stages[i].nir, nir_shader_get_entrypoint(stages[i].nir)); diff --git a/src/amd/vulkan/radv_pipeline_rt.c b/src/amd/vulkan/radv_pipeline_rt.c index f28769cedc8..cc148bda580 100644 --- a/src/amd/vulkan/radv_pipeline_rt.c +++ b/src/amd/vulkan/radv_pipeline_rt.c @@ -415,7 +415,7 @@ radv_rt_nir_to_asm(struct radv_device *device, struct vk_pipeline_cache *cache, temp_stage.nir = shaders[i]; radv_nir_lower_rt_abi(temp_stage.nir, pCreateInfo, &temp_stage.args, &stage->info, stack_size, i > 0, device, pipeline, monolithic); - radv_optimize_nir(temp_stage.nir, pipeline_key->optimisations_disabled); + radv_optimize_nir(temp_stage.nir, stage->key.optimisations_disabled); radv_postprocess_nir(device, pipeline_key, &temp_stage); if (radv_can_dump_shader(device, temp_stage.nir, false)) @@ -558,7 +558,7 @@ radv_rt_compile_shaders(struct radv_device *device, struct vk_pipeline_cache *ca if (nir_needed) { rt_stages[idx].stack_size = stage->nir->scratch_size; rt_stages[idx].nir = radv_pipeline_cache_nir_to_handle(device, cache, stage->nir, rt_stages[idx].sha1, - !key->optimisations_disabled); + !stage->key.optimisations_disabled); } stage->feedback.duration += os_time_get_nano() - stage_start; diff --git a/src/amd/vulkan/radv_shader.c b/src/amd/vulkan/radv_shader.c index 0a17b643317..4564d6b4256 100644 --- a/src/amd/vulkan/radv_shader.c +++ b/src/amd/vulkan/radv_shader.c @@ -670,7 +670,7 @@ radv_shader_spirv_to_nir(struct radv_device *device, const struct radv_shader_st NIR_PASS(_, nir, nir_lower_load_const_to_scalar); NIR_PASS(_, nir, nir_opt_shrink_stores, !device->instance->drirc.disable_shrink_image_store); - if (!key->optimisations_disabled) + if (!stage->key.optimisations_disabled) radv_optimize_nir(nir, false); /* We call nir_lower_var_copies() after the first radv_optimize_nir() @@ -741,8 +741,8 @@ radv_shader_spirv_to_nir(struct radv_device *device, const struct radv_shader_st * bloat the instruction count of the loop and cause it to be * considered too large for unrolling. */ - if (ac_nir_lower_indirect_derefs(nir, device->physical_device->rad_info.gfx_level) && !key->optimisations_disabled && - nir->info.stage != MESA_SHADER_COMPUTE) { + if (ac_nir_lower_indirect_derefs(nir, device->physical_device->rad_info.gfx_level) && + !stage->key.optimisations_disabled && nir->info.stage != MESA_SHADER_COMPUTE) { /* Optimize the lowered code before the linking optimizations. */ radv_optimize_nir(nir, false); } @@ -2471,7 +2471,7 @@ radv_capture_shader_executable_info(struct radv_device *device, struct radv_shad static struct radv_shader_binary * shader_compile(struct radv_device *device, struct nir_shader *const *shaders, int shader_count, gl_shader_stage stage, const struct radv_shader_info *info, const struct radv_shader_args *args, - struct radv_nir_compiler_options *options) + const struct radv_shader_stage_key *stage_key, struct radv_nir_compiler_options *options) { struct radv_shader_debug_data debug_data = { .device = device, @@ -2494,7 +2494,7 @@ shader_compile(struct radv_device *device, struct nir_shader *const *shaders, in } else { struct aco_shader_info ac_info; struct aco_compiler_options ac_opts; - radv_aco_convert_opts(&ac_opts, options, args); + radv_aco_convert_opts(&ac_opts, options, args, stage_key); radv_aco_convert_shader_info(&ac_info, info, args, &device->cache_key, options->info->gfx_level); aco_compile_shader(&ac_opts, &ac_info, shader_count, shaders, &args->ac, &radv_aco_build_shader_binary, (void **)&binary); @@ -2524,7 +2524,7 @@ radv_shader_nir_to_asm(struct radv_device *device, struct radv_shader_stage *pl_ keep_shader_info, keep_statistic_info); struct radv_shader_binary *binary = - shader_compile(device, shaders, shader_count, stage, info, &pl_stage->args, &options); + shader_compile(device, shaders, shader_count, stage, info, &pl_stage->args, &pl_stage->key, &options); return binary; } @@ -2550,6 +2550,7 @@ struct radv_shader * radv_create_trap_handler_shader(struct radv_device *device) { gl_shader_stage stage = MESA_SHADER_COMPUTE; + struct radv_shader_stage_key stage_key = {0}; struct radv_shader_info info = {0}; struct radv_pipeline_key key = {0}; struct radv_nir_compiler_options options = {0}; @@ -2564,7 +2565,7 @@ radv_create_trap_handler_shader(struct radv_device *device) struct radv_shader_args args; radv_declare_shader_args(device, &key, &info, stage, MESA_SHADER_NONE, &args); - struct radv_shader_binary *binary = shader_compile(device, &b.shader, 1, stage, &info, &args, &options); + struct radv_shader_binary *binary = shader_compile(device, &b.shader, 1, stage, &info, &args, &stage_key, &options); struct radv_shader *shader; radv_shader_create_uncached(device, binary, false, NULL, &shader); @@ -2637,10 +2638,11 @@ radv_create_rt_prolog(struct radv_device *device) #endif struct radv_shader_binary *binary = NULL; + struct radv_shader_stage_key stage_key = {0}; struct aco_shader_info ac_info; struct aco_compiler_options ac_opts; radv_aco_convert_shader_info(&ac_info, &info, &in_args, &device->cache_key, options.info->gfx_level); - radv_aco_convert_opts(&ac_opts, &options, &in_args); + radv_aco_convert_opts(&ac_opts, &options, &in_args, &stage_key); aco_compile_rt_prolog(&ac_opts, &ac_info, &in_args.ac, &out_args.ac, &radv_aco_build_shader_binary, (void **)&binary); binary->info = info; @@ -2700,11 +2702,12 @@ radv_create_vs_prolog(struct radv_device *device, const struct radv_vs_prolog_ke #endif struct radv_shader_part_binary *binary = NULL; + struct radv_shader_stage_key stage_key = {0}; struct aco_shader_info ac_info; struct aco_vs_prolog_info ac_prolog_info; struct aco_compiler_options ac_opts; radv_aco_convert_shader_info(&ac_info, &info, &args, &device->cache_key, options.info->gfx_level); - radv_aco_convert_opts(&ac_opts, &options, &args); + radv_aco_convert_opts(&ac_opts, &options, &args, &stage_key); radv_aco_convert_vs_prolog_key(&ac_prolog_info, key, &args); aco_compile_vs_prolog(&ac_opts, &ac_info, &ac_prolog_info, &args.ac, &radv_aco_build_shader_part, (void **)&binary); @@ -2753,11 +2756,12 @@ radv_create_ps_epilog(struct radv_device *device, const struct radv_ps_epilog_ke #endif struct radv_shader_part_binary *binary = NULL; + struct radv_shader_stage_key stage_key = {0}; struct aco_shader_info ac_info; struct aco_ps_epilog_info ac_epilog_info = {0}; struct aco_compiler_options ac_opts; radv_aco_convert_shader_info(&ac_info, &info, &args, &device->cache_key, options.info->gfx_level); - radv_aco_convert_opts(&ac_opts, &options, &args); + radv_aco_convert_opts(&ac_opts, &options, &args, &stage_key); radv_aco_convert_ps_epilog_key(&ac_epilog_info, key, &args); aco_compile_ps_epilog(&ac_opts, &ac_info, &ac_epilog_info, &args.ac, &radv_aco_build_shader_part, (void **)&binary); @@ -2811,11 +2815,12 @@ radv_create_tcs_epilog(struct radv_device *device, const struct radv_tcs_epilog_ #endif struct radv_shader_part_binary *binary = NULL; + struct radv_shader_stage_key stage_key = {0}; struct aco_shader_info ac_info; struct aco_tcs_epilog_info ac_epilog_info; struct aco_compiler_options ac_opts; radv_aco_convert_shader_info(&ac_info, &info, &args, &device->cache_key, options.info->gfx_level); - radv_aco_convert_opts(&ac_opts, &options, &args); + radv_aco_convert_opts(&ac_opts, &options, &args, &stage_key); radv_aco_convert_tcs_epilog_key(&ac_epilog_info, key, &args); aco_compile_tcs_epilog(&ac_opts, &ac_info, &ac_epilog_info, &args.ac, &radv_aco_build_shader_part, (void **)&binary); diff --git a/src/amd/vulkan/radv_shader.h b/src/amd/vulkan/radv_shader.h index bfb49baf3f2..1ba3fd4e759 100644 --- a/src/amd/vulkan/radv_shader.h +++ b/src/amd/vulkan/radv_shader.h @@ -92,6 +92,8 @@ struct radv_shader_stage_key { uint8_t storage_robustness2 : 1; uint8_t uniform_robustness2 : 1; + + uint8_t optimisations_disabled : 1; }; struct radv_ps_epilog_key { @@ -115,7 +117,6 @@ struct radv_pipeline_key { uint32_t lib_flags : 4; /* VkGraphicsPipelineLibraryFlagBitsEXT */ uint32_t has_multiview_view_index : 1; - uint32_t optimisations_disabled : 1; uint32_t adjust_frag_coord_z : 1; uint32_t dynamic_patch_control_points : 1; uint32_t dynamic_rasterization_samples : 1;
