From 0692ae34e939845e5185d3bdd33ddfe4afcdb995 Mon Sep 17 00:00:00 2001 From: =?utf8?q?Marek=20Ol=C5=A1=C3=A1k?= Date: Thu, 12 Sep 2019 19:46:02 -0400 Subject: [PATCH] ac: move ac_get_num_physical_sgprs into radeon_info Reviewed-by: Timothy Arceri Reviewed-by: Samuel Pitoiset --- src/amd/common/ac_gpu_info.c | 11 +++++++++++ src/amd/common/ac_gpu_info.h | 13 +------------ src/amd/vulkan/radv_device.c | 2 +- src/amd/vulkan/radv_shader.c | 4 ++-- src/gallium/drivers/radeonsi/si_shader.c | 4 ++-- 5 files changed, 17 insertions(+), 17 deletions(-) diff --git a/src/amd/common/ac_gpu_info.c b/src/amd/common/ac_gpu_info.c index cc7cbc11ee6..9db6c330a1a 100644 --- a/src/amd/common/ac_gpu_info.c +++ b/src/amd/common/ac_gpu_info.c @@ -586,6 +586,17 @@ bool ac_query_gpu_info(int fd, void *dev_p, info->max_wave64_per_simd = info->family >= CHIP_POLARIS10 && info->family <= CHIP_VEGAM ? 8 : 10; + + /* The number is per SIMD. There is enough SGPRs for the maximum number + * of Wave32, which is double the number for Wave64. + */ + if (info->chip_class >= GFX10) + info->num_physical_sgprs_per_simd = 128 * info->max_wave64_per_simd * 2; + else if (info->chip_class >= GFX8) + info->num_physical_sgprs_per_simd = 800; + else + info->num_physical_sgprs_per_simd = 512; + return true; } diff --git a/src/amd/common/ac_gpu_info.h b/src/amd/common/ac_gpu_info.h index 7ab9bb11e49..680c588a67a 100644 --- a/src/amd/common/ac_gpu_info.h +++ b/src/amd/common/ac_gpu_info.h @@ -142,6 +142,7 @@ struct radeon_info { uint32_t max_se; /* shader engines */ uint32_t max_sh_per_se; /* shader arrays per shader engine */ uint32_t max_wave64_per_simd; + uint32_t num_physical_sgprs_per_simd; /* Render backends (color + depth blocks). */ uint32_t r300_num_gb_pipes; @@ -200,18 +201,6 @@ static inline unsigned ac_get_num_physical_vgprs(enum chip_class chip_class, return 256; } -static inline uint32_t -ac_get_num_physical_sgprs(const struct radeon_info *info) -{ - /* The number is per SIMD. There is enough SGPRs for the maximum number - * of Wave32, which is double the number for Wave64. - */ - if (info->chip_class >= GFX10) - return 128 * info->max_wave64_per_simd * 2; - - return info->chip_class >= GFX8 ? 800 : 512; -} - #ifdef __cplusplus } #endif diff --git a/src/amd/vulkan/radv_device.c b/src/amd/vulkan/radv_device.c index 6a36c9afa6b..567fe00ac0f 100644 --- a/src/amd/vulkan/radv_device.c +++ b/src/amd/vulkan/radv_device.c @@ -1292,7 +1292,7 @@ void radv_GetPhysicalDeviceProperties2( /* SGPR. */ properties->sgprsPerSimd = - ac_get_num_physical_sgprs(&pdevice->rad_info); + pdevice->rad_info.num_physical_sgprs_per_simd; properties->minSgprAllocation = pdevice->rad_info.chip_class >= GFX8 ? 16 : 8; properties->maxSgprAllocation = diff --git a/src/amd/vulkan/radv_shader.c b/src/amd/vulkan/radv_shader.c index b875b989088..98abe8cd437 100644 --- a/src/amd/vulkan/radv_shader.c +++ b/src/amd/vulkan/radv_shader.c @@ -1274,7 +1274,7 @@ radv_get_max_waves(struct radv_device *device, if (conf->num_sgprs) max_simd_waves = MIN2(max_simd_waves, - ac_get_num_physical_sgprs(&device->physical_device->rad_info) / + device->physical_device->rad_info.num_physical_sgprs_per_simd / conf->num_sgprs); if (conf->num_vgprs) @@ -1372,7 +1372,7 @@ radv_GetShaderInfoAMD(VkDevice _device, VkShaderStatisticsInfoAMD statistics = {}; statistics.shaderStageMask = shaderStage; statistics.numPhysicalVgprs = RADV_NUM_PHYSICAL_VGPRS; - statistics.numPhysicalSgprs = ac_get_num_physical_sgprs(&device->physical_device->rad_info); + statistics.numPhysicalSgprs = device->physical_device->rad_info.num_physical_sgprs_per_simd; statistics.numAvailableSgprs = statistics.numPhysicalSgprs; if (stage == MESA_SHADER_COMPUTE) { diff --git a/src/gallium/drivers/radeonsi/si_shader.c b/src/gallium/drivers/radeonsi/si_shader.c index de1ad16a34b..fde6801fdce 100644 --- a/src/gallium/drivers/radeonsi/si_shader.c +++ b/src/gallium/drivers/radeonsi/si_shader.c @@ -5454,7 +5454,7 @@ static void si_calculate_max_simd_waves(struct si_shader *shader) if (conf->num_sgprs) { max_simd_waves = MIN2(max_simd_waves, - ac_get_num_physical_sgprs(&sscreen->info) / conf->num_sgprs); + sscreen->info.num_physical_sgprs_per_simd / conf->num_sgprs); } if (conf->num_vgprs) { @@ -7178,7 +7178,7 @@ int si_compile_tgsi_shader(struct si_screen *sscreen, unsigned wave_size = sscreen->compute_wave_size; unsigned max_vgprs = ac_get_num_physical_vgprs(sscreen->info.chip_class, wave_size); - unsigned max_sgprs = ac_get_num_physical_sgprs(&sscreen->info); + unsigned max_sgprs = sscreen->info.num_physical_sgprs_per_simd; unsigned max_sgprs_per_wave = 128; unsigned simds_per_tg = 4; /* assuming WGP mode on gfx10 */ unsigned threads_per_tg = si_get_max_workgroup_size(shader); -- 2.30.2