libagx: don't use get_group_id()

We don't want to see base_wg intrinsics.

Fixes: d22f936019 ("nir: remove workgroup_id_zero_base")
Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/29179>

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

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;