Mesa (master): radv: round vgprs/sgprs before calculating max_waves
GitLab Mirror
gitlab-mirror at kemper.freedesktop.org
Wed Oct 23 18:28:03 UTC 2019
Module: Mesa
Branch: master
Commit: 7453c1adff9d8a9e09cd7585e05c4db1c70870be
URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=7453c1adff9d8a9e09cd7585e05c4db1c70870be
Author: Rhys Perry <pendingchaos02 at gmail.com>
Date: Fri Oct 18 21:13:44 2019 +0100
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 <pendingchaos02 at gmail.com>
Reviewed-by: Daniel Schürmann <daniel at schuermann.dev>
---
src/amd/vulkan/radv_shader.c | 12 ++++++++----
1 file 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.
More information about the mesa-commit
mailing list