diff options
| author | Rhys Perry <[email protected]> | 2019-10-18 21:13:44 +0100 |
|---|---|---|
| committer | Rhys Perry <[email protected]> | 2019-10-23 19:11:20 +0100 |
| commit | 7453c1adff9d8a9e09cd7585e05c4db1c70870be (patch) | |
| tree | 5230bbad7ed4093ea1ed7d1489eedca4dec45316 /src/amd/vulkan/radv_shader.c | |
| parent | 254d9976b680c02822a7b6781589d92dbf7f55e4 (diff) | |
radv: round vgprs/sgprs before calculating max_waves
Note that ACO doesn't correctly round SGPR counts on GFX8/GFX9.
pipeline-db (ACO/Vega):
SGPRS: 11000 -> 11000 (0.00 %)
VGPRS: 3120 -> 3120 (0.00 %)
Spilled SGPRs: 0 -> 0 (0.00 %)
Spilled VGPRs: 0 -> 0 (0.00 %)
Private memory VGPRs: 0 -> 0 (0.00 %)
Scratch size: 0 -> 0 (0.00 %) dwords per thread
Code Size: 164328 -> 164328 (0.00 %) bytes
LDS: 0 -> 0 (0.00 %) blocks
Max Waves: 1125 -> 1000 (-11.11 %)
v2: consider wave32
Signed-off-by: Rhys Perry <[email protected]>
Reviewed-by: Daniel Schürmann <[email protected]>
Diffstat (limited to 'src/amd/vulkan/radv_shader.c')
| -rw-r--r-- | src/amd/vulkan/radv_shader.c | 12 |
1 files changed, 8 insertions, 4 deletions
diff --git a/src/amd/vulkan/radv_shader.c b/src/amd/vulkan/radv_shader.c index 5f962ed160b..06e8edf9ea2 100644 --- a/src/amd/vulkan/radv_shader.c +++ b/src/amd/vulkan/radv_shader.c @@ -1289,16 +1289,20 @@ radv_get_max_waves(struct radv_device *device, DIV_ROUND_UP(max_workgroup_size, wave_size); } - if (conf->num_sgprs) + if (conf->num_sgprs) { + 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 / - conf->num_sgprs); + sgprs); + } - if (conf->num_vgprs) + if (conf->num_vgprs) { + unsigned vgprs = align(conf->num_vgprs, wave_size == 32 ? 8 : 4); max_simd_waves = MIN2(max_simd_waves, - RADV_NUM_PHYSICAL_VGPRS / conf->num_vgprs); + RADV_NUM_PHYSICAL_VGPRS / vgprs); + } /* LDS is 64KB per CU (4 SIMDs), divided into 16KB blocks per SIMD * that PS can use. |