AlexIndustrial/mesa
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

agx: lower nir_intrinsic_load_num_subgroups

Browse Source 
Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/29179>
This commit is contained in:
Alyssa Rosenzweig
2024-04-19 22:03:37 -04:00
committed by Marge Bot
parent 765bf9a25b
commit 3efa723b63
1 changed files with 20 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/asahi/compiler/agx_nir_lower_subgroups.c
+20
View File
@@ -9,6 +9,13 @@
#include "nir_builder_opcodes.h"
#include "nir_intrinsics.h"
/* XXX: cribbed from nak, move to common */
static nir_def *
nir_udiv_round_up(nir_builder *b, nir_def *n, nir_def *d)
{
return nir_udiv(b, nir_iadd(b, n, nir_iadd_imm(b, d, -1)), d);
}
static bool
lower(nir_builder *b, nir_intrinsic_instr *intr, void *data)
{
@@ -67,6 +74,19 @@ lower(nir_builder *b, nir_intrinsic_instr *intr, void *data)
return true;
}
case nir_intrinsic_load_num_subgroups: {
nir_def *workgroup_size = nir_load_workgroup_size(b);
workgroup_size = nir_imul(b,
nir_imul(b, nir_channel(b, workgroup_size, 0),
nir_channel(b, workgroup_size, 1)),
nir_channel(b, workgroup_size, 2));
nir_def *subgroup_size = nir_imm_int(b, 32);
nir_def *num_subgroups =
nir_udiv_round_up(b, workgroup_size, subgroup_size);
nir_def_rewrite_uses(&intr->def, num_subgroups);
return true;
}
default:
return false;
}