From df61444ac47244f60db99f16d4e7c1e6ad403a4b Mon Sep 17 00:00:00 2001 From: Rhys Perry Date: Thu, 28 Jan 2021 11:42:35 +0000 Subject: [PATCH] radv: switch MaxWaves statistic to wave32 waves Always return the wave32 waves instead of wave64 waves because the wave32 wave count is more precise in the case of wave32. This also fixes usage of lds_per_wave in wave32. Signed-off-by: Rhys Perry Reviewed-by: Samuel Pitoiset Part-of: --- src/amd/vulkan/radv_shader.c | 13 +++++++------ 1 file changed, 7 insertions(+), 6 deletions(-) diff --git a/src/amd/vulkan/radv_shader.c b/src/amd/vulkan/radv_shader.c index 46a14b7..074d830 100644 --- a/src/amd/vulkan/radv_shader.c +++ b/src/amd/vulkan/radv_shader.c @@ -1569,7 +1569,8 @@ radv_get_max_waves(struct radv_device *device, unsigned max_simd_waves; unsigned lds_per_wave = 0; - max_simd_waves = device->physical_device->rad_info.max_wave64_per_simd; + max_simd_waves = device->physical_device->rad_info.max_wave64_per_simd * + (64 / wave_size); if (stage == MESA_SHADER_FRAGMENT) { lds_per_wave = conf->lds_size * lds_increment + @@ -1582,7 +1583,7 @@ radv_get_max_waves(struct radv_device *device, DIV_ROUND_UP(max_workgroup_size, wave_size); } - if (conf->num_sgprs) { + if (conf->num_sgprs && chip_class < GFX10) { unsigned sgprs = align(conf->num_sgprs, chip_class >= GFX8 ? 16 : 8); max_simd_waves = MIN2(max_simd_waves, @@ -1591,12 +1592,12 @@ radv_get_max_waves(struct radv_device *device, } if (conf->num_vgprs) { + unsigned physical_vgprs = device->physical_device->rad_info.num_physical_wave64_vgprs_per_simd * + (64 / wave_size); unsigned vgprs = align(conf->num_vgprs, wave_size == 32 ? 8 : 4); if (chip_class >= GFX10_3) vgprs = align(vgprs, wave_size == 32 ? 16 : 8); - max_simd_waves = - MIN2(max_simd_waves, - device->physical_device->rad_info.num_physical_wave64_vgprs_per_simd / vgprs); + max_simd_waves = MIN2(max_simd_waves, physical_vgprs / vgprs); } unsigned simd_per_workgroup = device->physical_device->rad_info.num_simd_per_compute_unit; @@ -1607,7 +1608,7 @@ radv_get_max_waves(struct radv_device *device, if (lds_per_wave) max_simd_waves = MIN2(max_simd_waves, max_lds_per_simd / lds_per_wave); - return max_simd_waves; + return chip_class >= GFX10 ? max_simd_waves * (wave_size / 32) : max_simd_waves; } VkResult -- 2.7.4