ac: move ac_get_num_physical_sgprs into radeon_info
Reviewed-by: Timothy Arceri <tarceri@itsqueeze.com> Reviewed-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
This commit is contained in:
@@ -586,6 +586,17 @@ bool ac_query_gpu_info(int fd, void *dev_p,
|
|||||||
|
|
||||||
info->max_wave64_per_simd = info->family >= CHIP_POLARIS10 &&
|
info->max_wave64_per_simd = info->family >= CHIP_POLARIS10 &&
|
||||||
info->family <= CHIP_VEGAM ? 8 : 10;
|
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;
|
return true;
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|||||||
@@ -142,6 +142,7 @@ struct radeon_info {
|
|||||||
uint32_t max_se; /* shader engines */
|
uint32_t max_se; /* shader engines */
|
||||||
uint32_t max_sh_per_se; /* shader arrays per shader engine */
|
uint32_t max_sh_per_se; /* shader arrays per shader engine */
|
||||||
uint32_t max_wave64_per_simd;
|
uint32_t max_wave64_per_simd;
|
||||||
|
uint32_t num_physical_sgprs_per_simd;
|
||||||
|
|
||||||
/* Render backends (color + depth blocks). */
|
/* Render backends (color + depth blocks). */
|
||||||
uint32_t r300_num_gb_pipes;
|
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;
|
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
|
#ifdef __cplusplus
|
||||||
}
|
}
|
||||||
#endif
|
#endif
|
||||||
|
|||||||
@@ -1292,7 +1292,7 @@ void radv_GetPhysicalDeviceProperties2(
|
|||||||
|
|
||||||
/* SGPR. */
|
/* SGPR. */
|
||||||
properties->sgprsPerSimd =
|
properties->sgprsPerSimd =
|
||||||
ac_get_num_physical_sgprs(&pdevice->rad_info);
|
pdevice->rad_info.num_physical_sgprs_per_simd;
|
||||||
properties->minSgprAllocation =
|
properties->minSgprAllocation =
|
||||||
pdevice->rad_info.chip_class >= GFX8 ? 16 : 8;
|
pdevice->rad_info.chip_class >= GFX8 ? 16 : 8;
|
||||||
properties->maxSgprAllocation =
|
properties->maxSgprAllocation =
|
||||||
|
|||||||
@@ -1274,7 +1274,7 @@ radv_get_max_waves(struct radv_device *device,
|
|||||||
if (conf->num_sgprs)
|
if (conf->num_sgprs)
|
||||||
max_simd_waves =
|
max_simd_waves =
|
||||||
MIN2(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);
|
conf->num_sgprs);
|
||||||
|
|
||||||
if (conf->num_vgprs)
|
if (conf->num_vgprs)
|
||||||
@@ -1372,7 +1372,7 @@ radv_GetShaderInfoAMD(VkDevice _device,
|
|||||||
VkShaderStatisticsInfoAMD statistics = {};
|
VkShaderStatisticsInfoAMD statistics = {};
|
||||||
statistics.shaderStageMask = shaderStage;
|
statistics.shaderStageMask = shaderStage;
|
||||||
statistics.numPhysicalVgprs = RADV_NUM_PHYSICAL_VGPRS;
|
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;
|
statistics.numAvailableSgprs = statistics.numPhysicalSgprs;
|
||||||
|
|
||||||
if (stage == MESA_SHADER_COMPUTE) {
|
if (stage == MESA_SHADER_COMPUTE) {
|
||||||
|
|||||||
@@ -5454,7 +5454,7 @@ static void si_calculate_max_simd_waves(struct si_shader *shader)
|
|||||||
if (conf->num_sgprs) {
|
if (conf->num_sgprs) {
|
||||||
max_simd_waves =
|
max_simd_waves =
|
||||||
MIN2(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) {
|
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 wave_size = sscreen->compute_wave_size;
|
||||||
unsigned max_vgprs = ac_get_num_physical_vgprs(sscreen->info.chip_class,
|
unsigned max_vgprs = ac_get_num_physical_vgprs(sscreen->info.chip_class,
|
||||||
wave_size);
|
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 max_sgprs_per_wave = 128;
|
||||||
unsigned simds_per_tg = 4; /* assuming WGP mode on gfx10 */
|
unsigned simds_per_tg = 4; /* assuming WGP mode on gfx10 */
|
||||||
unsigned threads_per_tg = si_get_max_workgroup_size(shader);
|
unsigned threads_per_tg = si_get_max_workgroup_size(shader);
|
||||||
|
|||||||
Reference in New Issue
Block a user