fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2025-12-25 10:50:10 +01:00
Code Issues Projects Releases Packages Wiki Activity Actions 1

radv: use lds_{encode,alloc}_granularity

Browse source 
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@gmail.com>
Reviewed-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8761>
This commit is contained in:
Rhys Perry 2021-01-28 12:03:03 +00:00
parent aa53335135
commit 267d7074d9
2 changed files with 18 additions and 20 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

36
src/amd/vulkan/radv_shader.c
View file

@ -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};

2
src/amd/vulkan/winsys/null/radv_null_winsys.c
View file

@ -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;