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;