diff --git a/src/asahi/libagx/compression.cl b/src/asahi/libagx/compression.cl index ca7f947ee74..e5e2fa673d0 100644 --- a/src/asahi/libagx/compression.cl +++ b/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; } } diff --git a/src/asahi/libagx/draws.cl b/src/asahi/libagx/draws.cl index 4663a229578..e2edfb64fa7 100644 --- a/src/asahi/libagx/draws.cl +++ b/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; diff --git a/src/asahi/libagx/geometry.cl b/src/asahi/libagx/geometry.cl index 46ccb076c87..9dfe6c4cb6f 100644 --- a/src/asahi/libagx/geometry.cl +++ b/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 */ diff --git a/src/asahi/libagx/query.cl b/src/asahi/libagx/query.cl index c1ec47d6b4c..f3647df5f61 100644 --- a/src/asahi/libagx/query.cl +++ b/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; } diff --git a/src/asahi/libagx/tessellator.cl b/src/asahi/libagx/tessellator.cl index 7d38107d481..a93c77e5bd1 100644 --- a/src/asahi/libagx/tessellator.cl +++ b/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);