From 87f2f77960b86079bc6f74fd59888f22a386bc87 Mon Sep 17 00:00:00 2001 From: Rhys Perry Date: Mon, 9 Dec 2024 14:37:49 +0000 Subject: [PATCH] aco: fix max_workgroup_count[0] This is necessary for radeonsi. fossil-db (navi21): Totals from 292 (0.37% of 79395) affected shaders: Instrs: 305965 -> 306182 (+0.07%); split: -0.00%, +0.07% CodeSize: 1624816 -> 1627212 (+0.15%); split: -0.00%, +0.15% Latency: 5244652 -> 5243587 (-0.02%); split: -0.07%, +0.05% InvThroughput: 1221089 -> 1225285 (+0.34%); split: -0.04%, +0.38% Copies: 22712 -> 22702 (-0.04%) PreSGPRs: 10713 -> 10712 (-0.01%) PreVGPRs: 10918 -> 10920 (+0.02%) VALU: 178613 -> 178836 (+0.12%) SALU: 43490 -> 43493 (+0.01%); split: -0.02%, +0.03% Signed-off-by: Rhys Perry Reviewed-by: Samuel Pitoiset Reviewed-by: Bas Nieuwenhuizen Part-of: --- src/amd/compiler/aco_instruction_selection_setup.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/amd/compiler/aco_instruction_selection_setup.cpp b/src/amd/compiler/aco_instruction_selection_setup.cpp index a322072194e..248e10548ea 100644 --- a/src/amd/compiler/aco_instruction_selection_setup.cpp +++ b/src/amd/compiler/aco_instruction_selection_setup.cpp @@ -357,7 +357,7 @@ init_context(isel_context* ctx, nir_shader* shader) ctx->ub_config.min_subgroup_size = ctx->program->wave_size; ctx->ub_config.max_subgroup_size = ctx->program->wave_size; ctx->ub_config.max_workgroup_invocations = 2048; - ctx->ub_config.max_workgroup_count[0] = 65535; + ctx->ub_config.max_workgroup_count[0] = 4294967295; ctx->ub_config.max_workgroup_count[1] = 65535; ctx->ub_config.max_workgroup_count[2] = 65535; ctx->ub_config.max_workgroup_size[0] = 2048;