Module: Mesa Branch: main Commit: 124b0039431be2bf0feeae8963a69a6965068338 URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=124b0039431be2bf0feeae8963a69a6965068338
Author: Samuel Pitoiset <[email protected]> Date: Fri Sep 24 15:10:32 2021 +0200 radv: store the CS subgroup size to radv_shader_info Signed-off-by: Samuel Pitoiset <[email protected]> Reviewed-by: Timur Kristóf <[email protected]> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13032> --- .../compiler/aco_instruction_selection_setup.cpp | 6 +- src/amd/vulkan/radv_pipeline.c | 76 +++++++++++----------- src/amd/vulkan/radv_shader.h | 7 +- 3 files changed, 44 insertions(+), 45 deletions(-) diff --git a/src/amd/compiler/aco_instruction_selection_setup.cpp b/src/amd/compiler/aco_instruction_selection_setup.cpp index 76e7c113d5a..8b7fe4f5395 100644 --- a/src/amd/compiler/aco_instruction_selection_setup.cpp +++ b/src/amd/compiler/aco_instruction_selection_setup.cpp @@ -468,9 +468,9 @@ init_context(isel_context* ctx, nir_shader* shader) ctx->range_ht = _mesa_pointer_hash_table_create(NULL); ctx->ub_config.min_subgroup_size = 64; ctx->ub_config.max_subgroup_size = 64; - if (ctx->shader->info.stage == MESA_SHADER_COMPUTE && ctx->options->key.cs.subgroup_size) { - ctx->ub_config.min_subgroup_size = ctx->options->key.cs.subgroup_size; - ctx->ub_config.max_subgroup_size = ctx->options->key.cs.subgroup_size; + if (ctx->shader->info.stage == MESA_SHADER_COMPUTE && ctx->args->shader_info->cs.subgroup_size) { + ctx->ub_config.min_subgroup_size = ctx->args->shader_info->cs.subgroup_size; + ctx->ub_config.max_subgroup_size = ctx->args->shader_info->cs.subgroup_size; } ctx->ub_config.max_workgroup_invocations = 2048; ctx->ub_config.max_workgroup_count[0] = 65535; diff --git a/src/amd/vulkan/radv_pipeline.c b/src/amd/vulkan/radv_pipeline.c index 857848425b8..e47bcc69b20 100644 --- a/src/amd/vulkan/radv_pipeline.c +++ b/src/amd/vulkan/radv_pipeline.c @@ -2819,45 +2819,16 @@ radv_fill_shader_keys(struct radv_device *device, struct radv_shader_variant_key keys[MESA_SHADER_FRAGMENT].fs.is_int10 = key->is_int10; keys[MESA_SHADER_FRAGMENT].fs.log2_ps_iter_samples = key->log2_ps_iter_samples; keys[MESA_SHADER_FRAGMENT].fs.num_samples = key->num_samples; - - if (nir[MESA_SHADER_COMPUTE]) { - unsigned subgroup_size = key->compute_subgroup_size; - unsigned req_subgroup_size = subgroup_size; - bool require_full_subgroups = key->require_full_subgroups; - - if (!subgroup_size) - subgroup_size = device->physical_device->cs_wave_size; - - unsigned local_size = nir[MESA_SHADER_COMPUTE]->info.workgroup_size[0] * - nir[MESA_SHADER_COMPUTE]->info.workgroup_size[1] * - nir[MESA_SHADER_COMPUTE]->info.workgroup_size[2]; - - /* Games don't always request full subgroups when they should, - * which can cause bugs if cswave32 is enabled. - */ - if (device->physical_device->cs_wave_size == 32 && - nir[MESA_SHADER_COMPUTE]->info.cs.uses_wide_subgroup_intrinsics && !req_subgroup_size && - local_size % RADV_SUBGROUP_SIZE == 0) - require_full_subgroups = true; - - if (require_full_subgroups && !req_subgroup_size) { - /* don't use wave32 pretending to be wave64 */ - subgroup_size = RADV_SUBGROUP_SIZE; - } - - keys[MESA_SHADER_COMPUTE].cs.subgroup_size = subgroup_size; - } } static uint8_t radv_get_wave_size(struct radv_device *device, const VkPipelineShaderStageCreateInfo *pStage, - gl_shader_stage stage, const struct radv_shader_variant_key *key, - const struct radv_shader_info *info) + gl_shader_stage stage, const struct radv_shader_info *info) { if (stage == MESA_SHADER_GEOMETRY && !info->is_ngg) return 64; else if (stage == MESA_SHADER_COMPUTE) { - return key->cs.subgroup_size; + return info->cs.subgroup_size; } else if (stage == MESA_SHADER_FRAGMENT) return device->physical_device->ps_wave_size; else @@ -2866,19 +2837,21 @@ radv_get_wave_size(struct radv_device *device, const VkPipelineShaderStageCreate static uint8_t radv_get_ballot_bit_size(struct radv_device *device, const VkPipelineShaderStageCreateInfo *pStage, - gl_shader_stage stage, const struct radv_shader_variant_key *key) + gl_shader_stage stage, const struct radv_shader_info *info) { - if (stage == MESA_SHADER_COMPUTE && key->cs.subgroup_size) - return key->cs.subgroup_size; + if (stage == MESA_SHADER_COMPUTE && info->cs.subgroup_size) + return info->cs.subgroup_size; return 64; } static void radv_fill_shader_info(struct radv_pipeline *pipeline, const VkPipelineShaderStageCreateInfo **pStages, + const struct radv_pipeline_key *pipeline_key, struct radv_shader_variant_key *keys, struct radv_shader_info *infos, nir_shader **nir) { + struct radv_device *device = pipeline->device; unsigned active_stages = 0; unsigned filled_stages = 0; @@ -2963,11 +2936,40 @@ radv_fill_shader_info(struct radv_pipeline *pipeline, radv_nir_shader_info_pass(pipeline->device, nir[i], pipeline->layout, &keys[i], &infos[i]); } + if (nir[MESA_SHADER_COMPUTE]) { + /* Variable workgroup size is not supported by Vulkan. */ + unsigned subgroup_size = pipeline_key->compute_subgroup_size; + unsigned req_subgroup_size = subgroup_size; + bool require_full_subgroups = pipeline_key->require_full_subgroups; + + if (!subgroup_size) + subgroup_size = device->physical_device->cs_wave_size; + + unsigned local_size = nir[MESA_SHADER_COMPUTE]->info.workgroup_size[0] * + nir[MESA_SHADER_COMPUTE]->info.workgroup_size[1] * + nir[MESA_SHADER_COMPUTE]->info.workgroup_size[2]; + + /* Games don't always request full subgroups when they should, + * which can cause bugs if cswave32 is enabled. + */ + if (device->physical_device->cs_wave_size == 32 && + nir[MESA_SHADER_COMPUTE]->info.cs.uses_wide_subgroup_intrinsics && !req_subgroup_size && + local_size % RADV_SUBGROUP_SIZE == 0) + require_full_subgroups = true; + + if (require_full_subgroups && !req_subgroup_size) { + /* don't use wave32 pretending to be wave64 */ + subgroup_size = RADV_SUBGROUP_SIZE; + } + + infos[MESA_SHADER_COMPUTE].cs.subgroup_size = subgroup_size; + } + for (int i = 0; i < MESA_SHADER_STAGES; i++) { if (nir[i]) { - infos[i].wave_size = radv_get_wave_size(pipeline->device, pStages[i], i, &keys[i], &infos[i]); + infos[i].wave_size = radv_get_wave_size(pipeline->device, pStages[i], i, &infos[i]); infos[i].ballot_bit_size = - radv_get_ballot_bit_size(pipeline->device, pStages[i], i, &keys[i]); + radv_get_ballot_bit_size(pipeline->device, pStages[i], i, &infos[i]); } } @@ -3444,7 +3446,7 @@ radv_create_shaders(struct radv_pipeline *pipeline, struct radv_device *device, } radv_fill_shader_keys(device, keys, pipeline_key, nir); - radv_fill_shader_info(pipeline, pStages, keys, infos, nir); + radv_fill_shader_info(pipeline, pStages, pipeline_key, keys, infos, nir); bool pipeline_has_ngg = (nir[MESA_SHADER_VERTEX] && keys[MESA_SHADER_VERTEX].vs_common_out.as_ngg) || (nir[MESA_SHADER_TESS_EVAL] && keys[MESA_SHADER_TESS_EVAL].vs_common_out.as_ngg); diff --git a/src/amd/vulkan/radv_shader.h b/src/amd/vulkan/radv_shader.h index b1ba250b1e5..ce9333c3975 100644 --- a/src/amd/vulkan/radv_shader.h +++ b/src/amd/vulkan/radv_shader.h @@ -100,17 +100,12 @@ struct radv_fs_variant_key { uint32_t is_int10; }; -struct radv_cs_variant_key { - uint8_t subgroup_size; -}; - struct radv_shader_variant_key { union { struct radv_vs_variant_key vs; struct radv_fs_variant_key fs; struct radv_tes_variant_key tes; struct radv_tcs_variant_key tcs; - struct radv_cs_variant_key cs; /* A common prefix of the vs and tes keys. */ struct radv_vs_out_key vs_common_out; @@ -346,6 +341,8 @@ struct radv_shader_info { bool uses_local_invocation_idx; unsigned block_size[3]; + uint8_t subgroup_size; + bool uses_sbt; bool uses_ray_launch_size; } cs;
