OpenHarmony/third_party_mesa3d
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

aco/gfx11+: use v_and_b32 to extract local id 0

Browse Source 
Foz-DB Navi31:
Totals from 2561 (3.23% of 79206) affected shaders:
CodeSize: 10399004 -> 10389120 (-0.10%)

Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/32532>
This commit is contained in:
Georg Lehmann
2024-12-06 20:44:47 +01:00
committed by Marge Bot
parent 1027b071f9
commit 4a977ea24f
1 changed files with 13 additions and 7 deletions
Download Patch File Download Diff File Expand all files Collapse all files

20
src/amd/compiler/aco_instruction_selection.cpp
View File

@@ -8232,13 +8232,19 @@ visit_intrinsic(isel_context* ctx, nir_intrinsic_instr* instr)
Temp local_ids[3];
/* Thread IDs are packed in VGPR0, 10 bits per component. */
for (uint32_t i = 0; i < 3; i++) {
if (i == 0 && ctx->shader->info.workgroup_size[1] == 1 &&
ctx->shader->info.workgroup_size[2] == 1 &&
!ctx->shader->info.workgroup_size_variable) {
local_ids[i] = get_arg(ctx, ctx->args->local_invocation_ids);
} else if (i == 2 || (i == 1 && ctx->shader->info.workgroup_size[2] == 1 &&
!ctx->shader->info.workgroup_size_variable)) {
local_ids[0] = get_arg(ctx, ctx->args->local_invocation_ids);
if (ctx->shader->info.workgroup_size[1] > 1 || ctx->shader->info.workgroup_size[2] > 1 ||
ctx->shader->info.workgroup_size_variable) {
unsigned size_x = ctx->shader->info.workgroup_size_variable
? 1024
: util_next_power_of_two(ctx->shader->info.workgroup_size[0]);
Temp mask = bld.copy(bld.def(s1), Operand::c32(size_x - 1));
local_ids[0] = bld.vop2(aco_opcode::v_and_b32, bld.def(v1), mask, local_ids[0]);
}
for (uint32_t i = 1; i < 3; i++) {
if (i == 2 || (i == 1 && ctx->shader->info.workgroup_size[2] == 1 &&
!ctx->shader->info.workgroup_size_variable)) {
local_ids[i] =
bld.vop2(aco_opcode::v_lshrrev_b32, bld.def(v1), Operand::c32(i * 10u),
get_arg(ctx, ctx->args->local_invocation_ids));