From d2cf17022da57f8363632307cc427878d1d4580d Mon Sep 17 00:00:00 2001 From: Alyssa Rosenzweig Date: Mon, 13 May 2024 14:23:21 -0400 Subject: [PATCH] libagx: don't use get_group_id() We don't want to see base_wg intrinsics. Fixes: d22f936019f ("nir: remove workgroup_id_zero_base") Signed-off-by: Alyssa Rosenzweig Part-of: --- src/asahi/lib/agx_nir_lower_gs.c | 3 ++- src/asahi/lib/shaders/geometry.cl | 3 +-- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/src/asahi/lib/agx_nir_lower_gs.c b/src/asahi/lib/agx_nir_lower_gs.c index 10093c57f22..047ffc2961f 100644 --- a/src/asahi/lib/agx_nir_lower_gs.c +++ b/src/asahi/lib/agx_nir_lower_gs.c @@ -1435,7 +1435,8 @@ agx_nir_prefix_sum_gs(nir_builder *b, const void *data) libagx_prefix_sum(b, load_geometry_param(b, count_buffer), load_geometry_param(b, input_primitives), - nir_imm_int(b, *words)); + nir_imm_int(b, *words), + nir_channel(b, nir_load_workgroup_id(b), 0)); } void diff --git a/src/asahi/lib/shaders/geometry.cl b/src/asahi/lib/shaders/geometry.cl index 2a58bd6e3d6..95edd21b0e5 100644 --- a/src/asahi/lib/shaders/geometry.cl +++ b/src/asahi/lib/shaders/geometry.cl @@ -544,11 +544,10 @@ libagx_work_group_scan_inclusive_add(uint x, local uint *scratch) } kernel void -libagx_prefix_sum(global uint *buffer, uint len, uint words) +libagx_prefix_sum(global uint *buffer, uint len, uint words, uint word) { local uint scratch[32]; uint tid = get_local_id(0); - uint word = get_group_id(0); /* Main loop: complete workgroups processing 1024 values at once */ uint i, count = 0;