Mesa (master): aco: fix num_waves on GFX10+
GitLab Mirror
gitlab-mirror at kemper.freedesktop.org
Wed Jan 20 17:00:53 UTC 2021
Module: Mesa
Branch: master
Commit: 489aa8c7cb7357b37f0e42771a9ca3a29eccb27b
URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=489aa8c7cb7357b37f0e42771a9ca3a29eccb27b
Author: Rhys Perry <pendingchaos02 at gmail.com>
Date: Tue Jan 19 11:37:52 2021 +0000
aco: fix num_waves on GFX10+
There are half the SIMDs per CU and physical_vgprs should be 512 instead
of 256.
fossil-db (GFX10.3):
Totals from 3622 (2.60% of 139391) affected shaders:
VGPRs: 298192 -> 289732 (-2.84%); split: -3.43%, +0.59%
CodeSize: 29443432 -> 29458388 (+0.05%); split: -0.00%, +0.06%
MaxWaves: 21703 -> 23395 (+7.80%); split: +7.84%, -0.05%
Instrs: 5677920 -> 5681438 (+0.06%); split: -0.01%, +0.07%
Cycles: 280715524 -> 280895676 (+0.06%); split: -0.00%, +0.07%
VMEM: 981142 -> 981894 (+0.08%); split: +0.18%, -0.10%
SMEM: 243315 -> 243454 (+0.06%); split: +0.07%, -0.02%
VClause: 88991 -> 89767 (+0.87%); split: -0.02%, +0.89%
SClause: 200660 -> 200659 (-0.00%); split: -0.00%, +0.00%
Copies: 430729 -> 434160 (+0.80%); split: -0.07%, +0.86%
Branches: 158004 -> 158021 (+0.01%); split: -0.01%, +0.02%
Signed-off-by: Rhys Perry <pendingchaos02 at gmail.com>
Reviewed-by: Daniel Schürmann <daniel at schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8523>
---
src/amd/compiler/aco_ir.cpp | 2 ++
src/amd/compiler/aco_ir.h | 1 +
src/amd/compiler/aco_live_var_analysis.cpp | 15 ++++++++-------
src/amd/compiler/aco_scheduler.cpp | 17 +++++++++++------
4 files changed, 22 insertions(+), 13 deletions(-)
diff --git a/src/amd/compiler/aco_ir.cpp b/src/amd/compiler/aco_ir.cpp
index 83ef05e7ebd..676a047c8b4 100644
--- a/src/amd/compiler/aco_ir.cpp
+++ b/src/amd/compiler/aco_ir.cpp
@@ -99,10 +99,12 @@ void init_program(Program *program, Stage stage, struct radv_shader_info *info,
program->has_16bank_lds = family == CHIP_KABINI || family == CHIP_STONEY;
program->vgpr_limit = 256;
+ program->physical_vgprs = 256;
program->vgpr_alloc_granule = 3;
if (chip_class >= GFX10) {
program->physical_sgprs = 2560; /* doesn't matter as long as it's at least 128 * 20 */
+ program->physical_vgprs = 512;
program->sgpr_alloc_granule = 127;
program->sgpr_limit = 106;
if (chip_class >= GFX10_3)
diff --git a/src/amd/compiler/aco_ir.h b/src/amd/compiler/aco_ir.h
index 69a9d977595..62b3f493fd6 100644
--- a/src/amd/compiler/aco_ir.h
+++ b/src/amd/compiler/aco_ir.h
@@ -1735,6 +1735,7 @@ public:
uint16_t vgpr_limit;
uint16_t sgpr_limit;
uint16_t physical_sgprs;
+ uint16_t physical_vgprs;
uint16_t sgpr_alloc_granule; /* minus one. must be power of two */
uint16_t vgpr_alloc_granule; /* minus one. must be power of two */
unsigned workgroup_size; /* if known; otherwise UINT_MAX */
diff --git a/src/amd/compiler/aco_live_var_analysis.cpp b/src/amd/compiler/aco_live_var_analysis.cpp
index bc713a1a188..1c041c9dbbb 100644
--- a/src/amd/compiler/aco_live_var_analysis.cpp
+++ b/src/amd/compiler/aco_live_var_analysis.cpp
@@ -314,7 +314,7 @@ uint16_t get_addr_sgpr_from_waves(Program *program, uint16_t max_waves)
uint16_t get_addr_vgpr_from_waves(Program *program, uint16_t max_waves)
{
- uint16_t vgprs = 256 / max_waves & ~program->vgpr_alloc_granule;
+ uint16_t vgprs = program->physical_vgprs / max_waves & ~program->vgpr_alloc_granule;
return std::min(vgprs, program->vgpr_limit);
}
@@ -325,7 +325,7 @@ void calc_min_waves(Program* program)
if (program->wave_size == 32)
waves_per_workgroup = DIV_ROUND_UP(waves_per_workgroup, 2);
- unsigned simd_per_cu = 4; /* TODO: different on Navi */
+ unsigned simd_per_cu = program->chip_class >= GFX10 ? 2 : 4;
bool wgp = program->chip_class >= GFX10; /* assume WGP is used on Navi */
unsigned simd_per_cu_wgp = wgp ? simd_per_cu * 2 : simd_per_cu;
@@ -334,11 +334,12 @@ void calc_min_waves(Program* program)
void update_vgpr_sgpr_demand(Program* program, const RegisterDemand new_demand)
{
- /* TODO: max_waves_per_simd, simd_per_cu and the number of physical vgprs for Navi */
- unsigned max_waves_per_simd = 10;
- if ((program->family >= CHIP_POLARIS10 && program->family <= CHIP_VEGAM) || program->chip_class >= GFX10_3)
+ unsigned max_waves_per_simd = program->chip_class == GFX10 ? 20 : 10;
+ if (program->chip_class >= GFX10_3)
+ max_waves_per_simd = 16;
+ else if (program->family >= CHIP_POLARIS10 && program->family <= CHIP_VEGAM)
max_waves_per_simd = 8;
- unsigned simd_per_cu = 4;
+ unsigned simd_per_cu = program->chip_class >= GFX10 ? 2 : 4;
bool wgp = program->chip_class >= GFX10; /* assume WGP is used on Navi */
unsigned simd_per_cu_wgp = wgp ? simd_per_cu * 2 : simd_per_cu;
@@ -350,7 +351,7 @@ void update_vgpr_sgpr_demand(Program* program, const RegisterDemand new_demand)
program->max_reg_demand = new_demand;
} else {
program->num_waves = program->physical_sgprs / get_sgpr_alloc(program, new_demand.sgpr);
- program->num_waves = std::min<uint16_t>(program->num_waves, 256 / get_vgpr_alloc(program, new_demand.vgpr));
+ program->num_waves = std::min<uint16_t>(program->num_waves, program->physical_vgprs / get_vgpr_alloc(program, new_demand.vgpr));
program->max_waves = max_waves_per_simd;
/* adjust max_waves for workgroup and LDS limits */
diff --git a/src/amd/compiler/aco_scheduler.cpp b/src/amd/compiler/aco_scheduler.cpp
index 0ae110f3209..fc9be0e0212 100644
--- a/src/amd/compiler/aco_scheduler.cpp
+++ b/src/amd/compiler/aco_scheduler.cpp
@@ -902,20 +902,25 @@ void schedule_program(Program *program, live& live_vars)
/* Allowing the scheduler to reduce the number of waves to as low as 5
* improves performance of Thrones of Britannia significantly and doesn't
* seem to hurt anything else. */
- if (program->num_waves <= 5)
+ //TODO: account for possible uneven num_waves on GFX10+
+ unsigned wave_fac = program->physical_vgprs / 256;
+ if (program->num_waves <= 5 * wave_fac)
ctx.num_waves = program->num_waves;
else if (demand.vgpr >= 29)
- ctx.num_waves = 5;
+ ctx.num_waves = 5 * wave_fac;
else if (demand.vgpr >= 25)
- ctx.num_waves = 6;
+ ctx.num_waves = 6 * wave_fac;
else
- ctx.num_waves = 7;
+ ctx.num_waves = 7 * wave_fac;
ctx.num_waves = std::max<uint16_t>(ctx.num_waves, program->min_waves);
ctx.num_waves = std::min<uint16_t>(ctx.num_waves, program->num_waves);
+ /* VMEM_MAX_MOVES and such assume pre-GFX10 wave count */
+ ctx.num_waves = std::max<uint16_t>(ctx.num_waves / wave_fac, 1);
+
assert(ctx.num_waves > 0);
- ctx.mv.max_registers = { int16_t(get_addr_vgpr_from_waves(program, ctx.num_waves) - 2),
- int16_t(get_addr_sgpr_from_waves(program, ctx.num_waves))};
+ ctx.mv.max_registers = { int16_t(get_addr_vgpr_from_waves(program, ctx.num_waves * wave_fac) - 2),
+ int16_t(get_addr_sgpr_from_waves(program, ctx.num_waves * wave_fac))};
for (Block& block : program->blocks)
schedule_block(ctx, program, &block, live_vars);
More information about the mesa-commit
mailing list