libagx: port to glsl-style compute builtins

Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Mary Guillemard <mary.guillemard@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/33099>

src/asahi/libagx/compression.cl

@@ -79,12 +79,10 @@ libagx_decompress(constant struct libagx_decompress_images *images,
                   uint16_t metadata_height_tl,
                   uint log2_samples__3 /* 1x, 2x, 4x */)
 {
-   uint3 coord_tl = (uint3)(get_group_id(0), get_group_id(1), get_group_id(2));
-   uint local_id = get_local_id(0);
    uint samples = 1 << log2_samples__3;
 
    /* Index into the metadata buffer */
-   uint index_tl = index_metadata(coord_tl, metadata_width_tl,
+   uint index_tl = index_metadata(cl_group_id, metadata_width_tl,
                                   metadata_height_tl, metadata_layer_stride_tl);
 
    /* If the tile is already uncompressed, there's nothing to do. */
@@ -92,12 +90,12 @@ libagx_decompress(constant struct libagx_decompress_images *images,
       return;
 
    /* Tiles are 16x16 */
-   uint2 coord_sa = (coord_tl.xy * 16);
-   uint layer = coord_tl.z;
+   uint2 coord_sa = (cl_group_id.xy * 16);
+   uint layer = cl_group_id.z;
 
    /* Since we use a 32x1 workgroup, each work-item handles half of a row. */
-   uint offs_y_sa = local_id >> 1;
-   uint offs_x_sa = (local_id & 1) ? 8 : 0;
+   uint offs_y_sa = cl_local_id.x >> 1;
+   uint offs_x_sa = (cl_local_id.x & 1) ? 8 : 0;
 
    int2 img_coord_sa_2d = convert_int2(coord_sa) + (int2)(offs_x_sa, offs_y_sa);
    int4 img_coord_sa = (int4)(img_coord_sa_2d.x, img_coord_sa_2d.y, layer, 0);
@@ -139,7 +137,7 @@ libagx_decompress(constant struct libagx_decompress_images *images,
    }
 
    /* We've replaced the body buffer. Mark the tile as uncompressed. */
-   if (local_id == 0) {
+   if (cl_local_id.x == 0) {
       metadata[index_tl] = tile_uncompressed;
    }
 }

src/asahi/libagx/draws.cl

@@ -16,7 +16,7 @@ libagx_predicate_indirect(global uint32_t *out, constant uint32_t *in,
                           constant uint32_t *draw_count, uint32_t stride_el,
                           uint indexed__2)
 {
-   uint draw = get_global_id(0);
+   uint draw = cl_global_id.x;
    uint words = indexed__2 ? 5 : 4;
    bool enabled = draw < *draw_count;
    out += draw * words;

src/asahi/libagx/geometry.cl

@@ -353,7 +353,7 @@ libagx_increment_ia_restart(global uint32_t *ia_vertices,
                             uint32_t restart_index, uint32_t index_size_B,
                             enum mesa_prim prim)
 {
-   uint tid = get_global_id(0);
+   uint tid = cl_global_id.x;
    unsigned count = draw[0];
    local uint scratch;
 
@@ -468,7 +468,7 @@ libagx_unroll_restart(global struct agx_geometry_state *heap,
 {
    uint32_t index_size_B = 1 << index_size_log2__3;
    enum mesa_prim mode = libagx_uncompact_prim(mode__11);
-   uint tid = get_local_id(0);
+   uint tid = cl_local_id.x;
    uint count = in_draw[0];
 
    local uintptr_t out_ptr;
@@ -716,7 +716,7 @@ KERNEL(1024)
 _libagx_prefix_sum(global uint *buffer, uint len, uint words, uint word)
 {
    local uint scratch[32];
-   uint tid = get_local_id(0);
+   uint tid = cl_local_id.x;
 
    /* Main loop: complete workgroups processing 1024 values at once */
    uint i, count = 0;
@@ -748,7 +748,7 @@ KERNEL(1024)
 libagx_prefix_sum_geom(constant struct agx_geometry_params *p)
 {
    _libagx_prefix_sum(p->count_buffer, p->input_primitives,
-                      p->count_buffer_stride / 4, get_group_id(0));
+                      p->count_buffer_stride / 4, cl_group_id.x);
 }
 
 KERNEL(1024)
@@ -760,7 +760,7 @@ libagx_prefix_sum_tess(global struct libagx_tess_args *p)
     * index buffer now. Elect a thread for the allocation.
     */
    barrier(CLK_LOCAL_MEM_FENCE);
-   if (get_local_id(0) != 0)
+   if (cl_local_id.x != 0)
       return;
 
    /* The last element of an inclusive prefix sum is the total sum */

src/asahi/libagx/query.cl

@@ -26,7 +26,7 @@ libagx_copy_query(global uint32_t *availability, global uint64_t *results,
                   uint16_t _64, uint16_t with_availability,
                   uint16_t reports_per_query)
 {
-   uint i = get_global_id(0);
+   uint i = cl_global_id.x;
    uint64_t dst = dst_addr + (((uint64_t)i) * dst_stride);
    uint32_t query = first_query + i;
 
@@ -88,7 +88,7 @@ libagx_copy_query_gl(global uint64_t *query, global uint64_t *dest,
 KERNEL(4)
 libagx_copy_xfb_counters(constant struct libagx_xfb_counter_copy *push)
 {
-   unsigned i = get_local_id(0);
+   unsigned i = cl_local_id.x;
 
    *(push->dest[i]) = push->src[i] ? *(push->src[i]) : 0;
 }
@@ -109,7 +109,7 @@ libagx_increment_cs_invocations(global uint *grid, global uint32_t *statistic,
 KERNEL(32)
 libagx_write_u32s(constant struct libagx_imm_write *p)
 {
-   uint id = get_global_id(0);
+   uint id = cl_global_id.x;
    *(p[id].address) = p[id].value;
 }
 

src/asahi/libagx/tessellator.cl

@@ -728,7 +728,7 @@ libagx_tess_isoline(constant struct libagx_tess_args *p,
                     enum libagx_tess_mode mode__2)
 {
    enum libagx_tess_mode mode = mode__2;
-   uint patch = get_global_id(0);
+   uint patch = cl_global_id.x;
    enum libagx_tess_partitioning partitioning = p->partitioning;
 
    bool lineDensityOdd;
@@ -830,7 +830,7 @@ libagx_tess_tri(constant struct libagx_tess_args *p,
                 enum libagx_tess_mode mode__2)
 {
    enum libagx_tess_mode mode = mode__2;
-   uint patch = get_global_id(0);
+   uint patch = cl_global_id.x;
    enum libagx_tess_partitioning partitioning = p->partitioning;
 
    global float *factors = tess_factors(p, patch);
@@ -1158,7 +1158,7 @@ libagx_tess_quad(constant struct libagx_tess_args *p,
                  enum libagx_tess_mode mode__2)
 {
    enum libagx_tess_mode mode = mode__2;
-   uint patch = get_global_id(0);
+   uint patch = cl_global_id.x;
    enum libagx_tess_partitioning partitioning = p->partitioning;
    global float *factors = tess_factors(p, patch);