radv: store the CS subgroup size to radv_shader_info
Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com> Reviewed-by: Timur Kristóf <timur.kristof@gmail.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13032>
This commit is contained in:
parent
c97147984b
commit
124b003943
|
@ -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;
|
||||
|
|
|
@ -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);
|
||||
|
|
|
@ -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;
|
||||
|
|
Loading…
Reference in New Issue