Mesa (master): radv: use lds_{encode,alloc}_granularity

[GitLab Mirror]

Module: Mesa
Branch: master
Commit: 267d7074d90c40d7488945ef11adb32618a9baf0
URL:    http://cgit.freedesktop.org/mesa/mesa/commit/?id=267d7074d90c40d7488945ef11adb32618a9baf0

Author: Rhys Perry <pendingchaos02 at gmail.com>
Date:   Thu Jan 28 12:03:03 2021 +0000

radv: use lds_{encode,alloc}_granularity

This fixes a issue in radv_get_max_waves() where it aligned the LDS
allocation to 512 bytes instead of 1024 on GFX10.3.

Signed-off-by: Rhys Perry <pendingchaos02 at gmail.com>
Reviewed-by: Samuel Pitoiset <samuel.pitoiset at gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8761>

---

 src/amd/vulkan/radv_shader.c                  | 36 ++++++++++++---------------
 src/amd/vulkan/winsys/null/radv_null_winsys.c |  2 ++
 2 files changed, 18 insertions(+), 20 deletions(-)

diff --git a/src/amd/vulkan/radv_shader.c b/src/amd/vulkan/radv_shader.c
index 074d8300579..c7a5ee94869 100644
--- a/src/amd/vulkan/radv_shader.c
+++ b/src/amd/vulkan/radv_shader.c
@@ -1209,8 +1209,8 @@ radv_shader_variant_create(struct radv_device *device,
 		}
 
 		if (rtld_binary.lds_size > 0) {
-			unsigned alloc_granularity = device->physical_device->rad_info.chip_class >= GFX7 ? 512 : 256;
-			config.lds_size = align(rtld_binary.lds_size, alloc_granularity) / alloc_granularity;
+			unsigned encode_granularity = device->physical_device->rad_info.lds_encode_granularity;
+			config.lds_size = align(rtld_binary.lds_size, encode_granularity) / encode_granularity;
 		}
 
 		variant->code_size = rtld_binary.rx_size;
@@ -1562,49 +1562,45 @@ radv_get_max_waves(struct radv_device *device,
                    struct radv_shader_variant *variant,
                    gl_shader_stage stage)
 {
-	enum chip_class chip_class = device->physical_device->rad_info.chip_class;
-	unsigned lds_increment = chip_class >= GFX7 ? 512 : 256;
+	struct radeon_info *info = &device->physical_device->rad_info;
+	enum chip_class chip_class = info->chip_class;
 	uint8_t wave_size = variant->info.wave_size;
 	struct ac_shader_config *conf = &variant->config;
 	unsigned max_simd_waves;
 	unsigned lds_per_wave = 0;
 
-	max_simd_waves = device->physical_device->rad_info.max_wave64_per_simd *
-			 (64 / wave_size);
+	max_simd_waves = info->max_wave64_per_simd * (64 / wave_size);
 
 	if (stage == MESA_SHADER_FRAGMENT) {
-		lds_per_wave = conf->lds_size * lds_increment +
-			       align(variant->info.ps.num_interp * 48,
-				     lds_increment);
+		lds_per_wave = conf->lds_size * info->lds_encode_granularity +
+			       variant->info.ps.num_interp * 48;
+		lds_per_wave = align(lds_per_wave, info->lds_alloc_granularity);
 	} else if (stage == MESA_SHADER_COMPUTE) {
 		unsigned max_workgroup_size =
 			radv_get_max_workgroup_size(chip_class, stage, variant->info.cs.block_size);
-		lds_per_wave = (conf->lds_size * lds_increment) /
-			       DIV_ROUND_UP(max_workgroup_size, wave_size);
+		lds_per_wave = align(conf->lds_size * info->lds_encode_granularity,
+				     info->lds_alloc_granularity);
+		lds_per_wave /= DIV_ROUND_UP(max_workgroup_size, wave_size);
 	}
 
 	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,
-			     device->physical_device->rad_info.num_physical_sgprs_per_simd /
-			     sgprs);
+		max_simd_waves = MIN2(max_simd_waves, info->num_physical_sgprs_per_simd / sgprs);
 	}
 
 	if (conf->num_vgprs) {
-		unsigned physical_vgprs = device->physical_device->rad_info.num_physical_wave64_vgprs_per_simd *
-					  (64 / wave_size);
+		unsigned physical_vgprs = 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, physical_vgprs / vgprs);
 	}
 
-	unsigned simd_per_workgroup = device->physical_device->rad_info.num_simd_per_compute_unit;
+	unsigned simd_per_workgroup = info->num_simd_per_compute_unit;
 	if (chip_class >= GFX10)
 		simd_per_workgroup *= 2; /* like lds_size_per_workgroup, assume WGP on GFX10+ */
 
-	unsigned max_lds_per_simd = device->physical_device->rad_info.lds_size_per_workgroup / simd_per_workgroup;
+	unsigned max_lds_per_simd = info->lds_size_per_workgroup / simd_per_workgroup;
 	if (lds_per_wave)
 		max_simd_waves = MIN2(max_simd_waves, max_lds_per_simd / lds_per_wave);
 
@@ -1635,7 +1631,7 @@ radv_GetShaderInfoAMD(VkDevice _device,
 		if (!pInfo) {
 			*pInfoSize = sizeof(VkShaderStatisticsInfoAMD);
 		} else {
-			unsigned lds_multiplier = device->physical_device->rad_info.chip_class >= GFX7 ? 512 : 256;
+			unsigned lds_multiplier = device->physical_device->rad_info.lds_encode_granularity;
 			struct ac_shader_config *conf = &variant->config;
 
 			VkShaderStatisticsInfoAMD statistics = {0};
diff --git a/src/amd/vulkan/winsys/null/radv_null_winsys.c b/src/amd/vulkan/winsys/null/radv_null_winsys.c
index 6320a975da9..99f324eebfb 100644
--- a/src/amd/vulkan/winsys/null/radv_null_winsys.c
+++ b/src/amd/vulkan/winsys/null/radv_null_winsys.c
@@ -128,6 +128,8 @@ static void radv_null_winsys_query_info(struct radeon_winsys *rws,
 	info->num_physical_wave64_vgprs_per_simd = info->chip_class >= GFX10 ? 512 : 256;
 	info->num_simd_per_compute_unit = info->chip_class >= GFX10 ? 2 : 4;
 	info->lds_size_per_workgroup = info->chip_class >= GFX10 ? 128 * 1024 : 64 * 1024;
+	info->lds_encode_granularity = info->chip_class >= GFX7 ? 128 * 4 : 64 * 4;
+	info->lds_alloc_granularity = info->chip_class >= GFX10_3 ? 256 * 4 : info->lds_encode_granularity;
 	info->max_render_backends = gpu_info[info->family].num_render_backends;
 
 	info->has_dedicated_vram = gpu_info[info->family].has_dedicated_vram;