xref: /aosp_15_r20/external/mesa3d/src/gallium/drivers/panfrost/pan_csf.c (revision 6104692788411f58d303aa86923a9ff6ecaded22)

/*
 * Copyright (C) 2023 Collabora Ltd.
 *
 * Permission is hereby granted, free of charge, to any person obtaining a
 * copy of this software and associated documentation files (the "Software"),
 * to deal in the Software without restriction, including without limitation
 * the rights to use, copy, modify, merge, publish, distribute, sublicense,
 * and/or sell copies of the Software, and to permit persons to whom the
 * Software is furnished to do so, subject to the following conditions:
 *
 * The above copyright notice and this permission notice (including the next
 * paragraph) shall be included in all copies or substantial portions of the
 * Software.
 *
 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL
 * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
 * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
 * SOFTWARE.
 */

#include "decode.h"

#include "drm-uapi/panthor_drm.h"

#include "genxml/cs_builder.h"
#include "panfrost/lib/genxml/cs_builder.h"

#include "pan_blitter.h"
#include "pan_cmdstream.h"
#include "pan_context.h"
#include "pan_csf.h"
#include "pan_job.h"

#if PAN_ARCH < 10
#error "CSF helpers are only used for gen >= 10"
#endif

static struct cs_buffer
csf_alloc_cs_buffer(void *cookie)
{
   assert(cookie && "Self-contained queues can't be extended.");

   struct panfrost_batch *batch = cookie;
   unsigned capacity = 4096;

   struct panfrost_ptr ptr =
      pan_pool_alloc_aligned(&batch->csf.cs_chunk_pool.base, capacity * 8, 64);

   return (struct cs_buffer){
      .cpu = ptr.cpu,
      .gpu = ptr.gpu,
      .capacity = capacity,
   };
}

void
GENX(csf_cleanup_batch)(struct panfrost_batch *batch)
{
   free(batch->csf.cs.builder);

   panfrost_pool_cleanup(&batch->csf.cs_chunk_pool);
}

void
GENX(csf_init_batch)(struct panfrost_batch *batch)
{
   struct panfrost_device *dev = pan_device(batch->ctx->base.screen);

   /* Initialize the CS chunk pool. */
   panfrost_pool_init(&batch->csf.cs_chunk_pool, NULL, dev, 0, 32768,
                      "CS chunk pool", false, true);

   /* Allocate and bind the command queue */
   struct cs_buffer queue = csf_alloc_cs_buffer(batch);
   const struct cs_builder_conf conf = {
      .nr_registers = 96,
      .nr_kernel_registers = 4,
      .alloc_buffer = csf_alloc_cs_buffer,
      .cookie = batch,
   };

   /* Setup the queue builder */
   batch->csf.cs.builder = malloc(sizeof(struct cs_builder));
   cs_builder_init(batch->csf.cs.builder, &conf, queue);
   cs_req_res(batch->csf.cs.builder,
              CS_COMPUTE_RES | CS_TILER_RES | CS_IDVS_RES | CS_FRAG_RES);

   /* Set up entries */
   struct cs_builder *b = batch->csf.cs.builder;
   cs_set_scoreboard_entry(b, 2, 0);

   batch->framebuffer = pan_pool_alloc_desc_aggregate(
      &batch->pool.base, PAN_DESC(FRAMEBUFFER), PAN_DESC(ZS_CRC_EXTENSION),
      PAN_DESC_ARRAY(MAX2(batch->key.nr_cbufs, 1), RENDER_TARGET));
   batch->tls = pan_pool_alloc_desc(&batch->pool.base, LOCAL_STORAGE);
}

static void
csf_prepare_qsubmit(struct panfrost_context *ctx,
                    struct drm_panthor_queue_submit *submit, uint8_t queue,
                    uint64_t cs_start, uint32_t cs_size,
                    struct drm_panthor_sync_op *syncs, uint32_t sync_count)
{
   struct panfrost_device *dev = pan_device(ctx->base.screen);

   *submit = (struct drm_panthor_queue_submit){
      .queue_index = queue,
      .stream_addr = cs_start,
      .stream_size = cs_size,
      .latest_flush = panthor_kmod_get_flush_id(dev->kmod.dev),
      .syncs = DRM_PANTHOR_OBJ_ARRAY(sync_count, syncs),
   };
}

static void
csf_prepare_gsubmit(struct panfrost_context *ctx,
                    struct drm_panthor_group_submit *gsubmit,
                    struct drm_panthor_queue_submit *qsubmits,
                    uint32_t qsubmit_count)
{
   *gsubmit = (struct drm_panthor_group_submit){
      .group_handle = ctx->csf.group_handle,
      .queue_submits = DRM_PANTHOR_OBJ_ARRAY(qsubmit_count, qsubmits),
   };
}

static int
csf_submit_gsubmit(struct panfrost_context *ctx,
                   struct drm_panthor_group_submit *gsubmit)
{
   struct panfrost_device *dev = pan_device(ctx->base.screen);
   int ret = 0;

   if (!ctx->is_noop) {
      ret = drmIoctl(panfrost_device_fd(dev), DRM_IOCTL_PANTHOR_GROUP_SUBMIT,
                     gsubmit);
   }

   if (ret)
      return errno;

   return 0;
}

static void
csf_emit_batch_end(struct panfrost_batch *batch)
{
   struct panfrost_device *dev = pan_device(batch->ctx->base.screen);
   struct cs_builder *b = batch->csf.cs.builder;

   /* Barrier to let everything finish */
   cs_wait_slots(b, BITFIELD_MASK(8), false);

   if (dev->debug & PAN_DBG_SYNC) {
      /* Get the CS state */
      batch->csf.cs.state = pan_pool_alloc_aligned(&batch->pool.base, 8, 8);
      memset(batch->csf.cs.state.cpu, ~0, 8);
      cs_move64_to(b, cs_reg64(b, 90), batch->csf.cs.state.gpu);
      cs_store_state(b, cs_reg64(b, 90), 0, MALI_CS_STATE_ERROR_STATUS,
                     cs_now());
   }

   /* Flush caches now that we're done (synchronous) */
   struct cs_index flush_id = cs_reg32(b, 74);
   cs_move32_to(b, flush_id, 0);
   cs_flush_caches(b, MALI_CS_FLUSH_MODE_CLEAN, MALI_CS_FLUSH_MODE_CLEAN, true,
                   flush_id, cs_defer(0, 0));
   cs_wait_slot(b, 0, false);

   /* Finish the command stream */
   assert(cs_is_valid(batch->csf.cs.builder));
   cs_finish(batch->csf.cs.builder);
}

static int
csf_submit_collect_wait_ops(struct panfrost_batch *batch,
                            struct util_dynarray *syncops,
                            uint32_t vm_sync_handle)
{
   struct panfrost_context *ctx = batch->ctx;
   struct panfrost_device *dev = pan_device(ctx->base.screen);
   uint64_t vm_sync_wait_point = 0, bo_sync_point;
   uint32_t bo_sync_handle;
   int ret;

   /* We don't wait on BOs attached to the various batch pools, because those
    * are private to the batch, and are guaranteed to be idle at allocation
    * time. We need to iterate over other BOs accessed by the batch though,
    * to add the corresponding wait operations.
    */
   util_dynarray_foreach(&batch->bos, pan_bo_access, ptr) {
      unsigned i = ptr - util_dynarray_element(&batch->bos, pan_bo_access, 0);
      pan_bo_access flags = *ptr;

      if (!flags)
         continue;

      /* Update the BO access flags so that panfrost_bo_wait() knows
       * about all pending accesses.
       * We only keep the READ/WRITE info since this is all the BO
       * wait logic cares about.
       * We also preserve existing flags as this batch might not
       * be the first one to access the BO.
       */
      struct panfrost_bo *bo = pan_lookup_bo(dev, i);

      ret = panthor_kmod_bo_get_sync_point(bo->kmod_bo, &bo_sync_handle,
                                           &bo_sync_point,
                                           !(flags & PAN_BO_ACCESS_WRITE));
      if (ret)
         return ret;

      if (bo_sync_handle == vm_sync_handle) {
         vm_sync_wait_point = MAX2(vm_sync_wait_point, bo_sync_point);
         continue;
      }

      assert(bo_sync_point == 0 || !bo->kmod_bo->exclusive_vm);

      struct drm_panthor_sync_op waitop = {
         .flags =
            DRM_PANTHOR_SYNC_OP_WAIT |
            (bo_sync_point ? DRM_PANTHOR_SYNC_OP_HANDLE_TYPE_TIMELINE_SYNCOBJ
                           : DRM_PANTHOR_SYNC_OP_HANDLE_TYPE_SYNCOBJ),
         .handle = bo_sync_handle,
         .timeline_value = bo_sync_point,
      };

      util_dynarray_append(syncops, struct drm_panthor_sync_op, waitop);
   }

   if (vm_sync_wait_point > 0) {
      struct drm_panthor_sync_op waitop = {
         .flags = DRM_PANTHOR_SYNC_OP_WAIT |
                  DRM_PANTHOR_SYNC_OP_HANDLE_TYPE_TIMELINE_SYNCOBJ,
         .handle = vm_sync_handle,
         .timeline_value = vm_sync_wait_point,
      };

      util_dynarray_append(syncops, struct drm_panthor_sync_op, waitop);
   }

   if (ctx->in_sync_fd >= 0) {
      ret = drmSyncobjImportSyncFile(panfrost_device_fd(dev), ctx->in_sync_obj,
                                     ctx->in_sync_fd);
      if (ret)
         return ret;

      struct drm_panthor_sync_op waitop = {
         .flags =
            DRM_PANTHOR_SYNC_OP_WAIT | DRM_PANTHOR_SYNC_OP_HANDLE_TYPE_SYNCOBJ,
         .handle = ctx->in_sync_obj,
      };

      util_dynarray_append(syncops, struct drm_panthor_sync_op, waitop);

      close(ctx->in_sync_fd);
      ctx->in_sync_fd = -1;
   }

   return 0;
}

static int
csf_attach_sync_points(struct panfrost_batch *batch, uint32_t vm_sync_handle,
                       uint64_t vm_sync_signal_point)
{
   struct panfrost_context *ctx = batch->ctx;
   struct panfrost_device *dev = pan_device(ctx->base.screen);
   int ret;

   /* There should be no invisble allocation on CSF. */
   assert(batch->invisible_pool.bos.size == 0);

   /* Attach sync points to batch-private BOs first. We assume BOs can
    * be written by the GPU to keep things simple.
    */
   util_dynarray_foreach(&batch->pool.bos, struct panfrost_bo *, bo) {
      (*bo)->gpu_access |= PAN_BO_ACCESS_RW;
      ret = panthor_kmod_bo_attach_sync_point((*bo)->kmod_bo, vm_sync_handle,
                                              vm_sync_signal_point, true);
      if (ret)
         return ret;
   }

   util_dynarray_foreach(&batch->csf.cs_chunk_pool.bos, struct panfrost_bo *,
                         bo) {
      (*bo)->gpu_access |= PAN_BO_ACCESS_RW;
      ret = panthor_kmod_bo_attach_sync_point((*bo)->kmod_bo, vm_sync_handle,
                                              vm_sync_signal_point, true);
      if (ret)
         return ret;
   }

   /* Attach the VM sync point to all resources accessed by the batch. */
   util_dynarray_foreach(&batch->bos, pan_bo_access, ptr) {
      unsigned i = ptr - util_dynarray_element(&batch->bos, pan_bo_access, 0);
      pan_bo_access flags = *ptr;

      if (!flags)
         continue;

      struct panfrost_bo *bo = pan_lookup_bo(dev, i);

      bo->gpu_access |= flags & (PAN_BO_ACCESS_RW);
      ret = panthor_kmod_bo_attach_sync_point(bo->kmod_bo, vm_sync_handle,
                                              vm_sync_signal_point,
                                              flags & PAN_BO_ACCESS_WRITE);
      if (ret)
         return ret;
   }

   /* And finally transfer the VM sync point to the context syncobj. */
   return drmSyncobjTransfer(panfrost_device_fd(dev), ctx->syncobj, 0,
                             vm_sync_handle, vm_sync_signal_point, 0);
}

static void
csf_check_ctx_state_and_reinit(struct panfrost_context *ctx)
{
   struct panfrost_device *dev = pan_device(ctx->base.screen);
   struct drm_panthor_group_get_state state = {
      .group_handle = ctx->csf.group_handle,
   };
   int ret;

   ret = drmIoctl(panfrost_device_fd(dev), DRM_IOCTL_PANTHOR_GROUP_GET_STATE,
                  &state);
   if (ret) {
      mesa_loge("DRM_IOCTL_PANTHOR_GROUP_GET_STATE failed (err=%d)", errno);
      return;
   }

   /* Context is still usable. This was a transient error. */
   if (state.state == 0)
      return;

   /* If the VM is unusable, we can't do much, as this is shared between all
    * contexts, and restoring the VM state is non-trivial.
    */
   if (pan_kmod_vm_query_state(dev->kmod.vm) != PAN_KMOD_VM_USABLE) {
      mesa_loge("VM became unusable, we can't reset the context");
      assert(!"VM became unusable, we can't reset the context");
   }

   panfrost_context_reinit(ctx);
}

static void
csf_submit_wait_and_dump(struct panfrost_batch *batch,
                         const struct drm_panthor_group_submit *gsubmit,
                         uint32_t vm_sync_handle, uint64_t vm_sync_signal_point)
{
   struct panfrost_context *ctx = batch->ctx;
   struct panfrost_device *dev = pan_device(ctx->base.screen);
   bool wait = (dev->debug & (PAN_DBG_TRACE | PAN_DBG_SYNC)) && !ctx->is_noop;
   bool dump = (dev->debug & PAN_DBG_TRACE);
   bool crash = false;

   if (!wait && !dump)
      return;

   /* Wait so we can get errors reported back */
   if (wait) {
      int ret =
         drmSyncobjTimelineWait(panfrost_device_fd(dev), &vm_sync_handle,
                                &vm_sync_signal_point, 1, INT64_MAX, 0, NULL);
      assert(ret >= 0);
   }

   /* Jobs won't be complete if blackhole rendering, that's ok */
   if (!ctx->is_noop && (dev->debug & PAN_DBG_SYNC) &&
       *((uint64_t *)batch->csf.cs.state.cpu) != 0) {
      crash = true;
      dump = true;
   }

   if (dump) {
      const struct drm_panthor_queue_submit *qsubmits =
         (void *)(uintptr_t)gsubmit->queue_submits.array;

      for (unsigned i = 0; i < gsubmit->queue_submits.count; i++) {
         uint32_t regs[256] = {0};
         pandecode_cs(dev->decode_ctx, qsubmits[i].stream_addr,
                      qsubmits[i].stream_size, panfrost_device_gpu_id(dev),
                      regs);
      }

      if (dev->debug & PAN_DBG_DUMP)
         pandecode_dump_mappings(dev->decode_ctx);
   }

   if (crash) {
      fprintf(stderr, "Incomplete job or timeout\n");
      fflush(NULL);
      abort();
   }
}

int
GENX(csf_submit_batch)(struct panfrost_batch *batch)
{
   /* Close the batch before submitting. */
   csf_emit_batch_end(batch);

   uint64_t cs_start = cs_root_chunk_gpu_addr(batch->csf.cs.builder);
   uint32_t cs_size = cs_root_chunk_size(batch->csf.cs.builder);
   struct panfrost_context *ctx = batch->ctx;
   struct panfrost_device *dev = pan_device(ctx->base.screen);
   uint32_t vm_sync_handle = panthor_kmod_vm_sync_handle(dev->kmod.vm);
   struct util_dynarray syncops;
   int ret;

   util_dynarray_init(&syncops, NULL);

   ret = csf_submit_collect_wait_ops(batch, &syncops, vm_sync_handle);
   if (ret)
      goto out_free_syncops;

   uint64_t vm_sync_cur_point = panthor_kmod_vm_sync_lock(dev->kmod.vm);
   uint64_t vm_sync_signal_point = vm_sync_cur_point + 1;

   struct drm_panthor_sync_op signalop = {
      .flags = DRM_PANTHOR_SYNC_OP_SIGNAL |
               DRM_PANTHOR_SYNC_OP_HANDLE_TYPE_TIMELINE_SYNCOBJ,
      .handle = vm_sync_handle,
      .timeline_value = vm_sync_signal_point,
   };

   util_dynarray_append(&syncops, struct drm_panthor_sync_op, signalop);

   struct drm_panthor_queue_submit qsubmit;
   struct drm_panthor_group_submit gsubmit;

   csf_prepare_qsubmit(
      ctx, &qsubmit, 0, cs_start, cs_size, util_dynarray_begin(&syncops),
      util_dynarray_num_elements(&syncops, struct drm_panthor_sync_op));
   csf_prepare_gsubmit(ctx, &gsubmit, &qsubmit, 1);
   ret = csf_submit_gsubmit(ctx, &gsubmit);
   panthor_kmod_vm_sync_unlock(dev->kmod.vm,
                               ret ? vm_sync_cur_point : vm_sync_signal_point);

   if (!ret) {
      csf_submit_wait_and_dump(batch, &gsubmit, vm_sync_handle,
                               vm_sync_signal_point);
      ret = csf_attach_sync_points(batch, vm_sync_handle, vm_sync_signal_point);
   } else {
      csf_check_ctx_state_and_reinit(batch->ctx);
   }

out_free_syncops:
   util_dynarray_fini(&syncops);
   return ret;
}

void
GENX(csf_preload_fb)(struct panfrost_batch *batch, struct pan_fb_info *fb)
{
   struct panfrost_device *dev = pan_device(batch->ctx->base.screen);

   GENX(pan_preload_fb)
   (&dev->blitter, &batch->pool.base, fb, 0, batch->tls.gpu, NULL);
}

void
GENX(csf_emit_fragment_job)(struct panfrost_batch *batch,
                            const struct pan_fb_info *pfb)
{
   struct cs_builder *b = batch->csf.cs.builder;

   if (batch->draw_count > 0) {
      /* Finish tiling and wait for IDVS and tiling */
      cs_finish_tiling(b, false);
      cs_wait_slot(b, 2, false);
      cs_vt_end(b, cs_now());
   }

   /* Set up the fragment job */
   cs_move64_to(b, cs_reg64(b, 40), batch->framebuffer.gpu);
   cs_move32_to(b, cs_reg32(b, 42), (batch->miny << 16) | batch->minx);
   cs_move32_to(b, cs_reg32(b, 43),
                ((batch->maxy - 1) << 16) | (batch->maxx - 1));

   /* Run the fragment job and wait */
   cs_run_fragment(b, false, MALI_TILE_RENDER_ORDER_Z_ORDER, false);
   cs_wait_slot(b, 2, false);

   /* Gather freed heap chunks and add them to the heap context free list
    * so they can be re-used next time the tiler heap runs out of chunks.
    * That's what cs_finish_fragment() is all about. The list of freed
    * chunks is in the tiler context descriptor
    * (completed_{top,bottom fields}). */
   if (batch->draw_count > 0) {
      assert(batch->tiler_ctx.valhall.desc);
      cs_move64_to(b, cs_reg64(b, 90), batch->tiler_ctx.valhall.desc);
      cs_load_to(b, cs_reg_tuple(b, 86, 4), cs_reg64(b, 90), BITFIELD_MASK(4),
                 40);
      cs_wait_slot(b, 0, false);
      cs_finish_fragment(b, true, cs_reg64(b, 86), cs_reg64(b, 88), cs_now());
   }
}

static void
csf_emit_shader_regs(struct panfrost_batch *batch, enum pipe_shader_type stage,
                     mali_ptr shader)
{
   mali_ptr resources = panfrost_emit_resources(batch, stage);

   assert(stage == PIPE_SHADER_VERTEX || stage == PIPE_SHADER_FRAGMENT ||
          stage == PIPE_SHADER_COMPUTE);

   unsigned offset = (stage == PIPE_SHADER_FRAGMENT) ? 4 : 0;
   unsigned fau_count = DIV_ROUND_UP(batch->nr_push_uniforms[stage], 2);

   struct cs_builder *b = batch->csf.cs.builder;
   cs_move64_to(b, cs_reg64(b, 0 + offset), resources);
   cs_move64_to(b, cs_reg64(b, 8 + offset),
                batch->push_uniforms[stage] | ((uint64_t)fau_count << 56));
   cs_move64_to(b, cs_reg64(b, 16 + offset), shader);
}

void
GENX(csf_launch_grid)(struct panfrost_batch *batch,
                      const struct pipe_grid_info *info)
{
   /* Empty compute programs are invalid and don't make sense */
   if (batch->rsd[PIPE_SHADER_COMPUTE] == 0)
      return;

   struct panfrost_context *ctx = batch->ctx;
   struct panfrost_device *dev = pan_device(ctx->base.screen);
   struct panfrost_compiled_shader *cs = ctx->prog[PIPE_SHADER_COMPUTE];
   struct cs_builder *b = batch->csf.cs.builder;

   csf_emit_shader_regs(batch, PIPE_SHADER_COMPUTE,
                        batch->rsd[PIPE_SHADER_COMPUTE]);

   cs_move64_to(b, cs_reg64(b, 24), batch->tls.gpu);

   /* Global attribute offset */
   cs_move32_to(b, cs_reg32(b, 32), 0);

   /* Compute workgroup size */
   uint32_t wg_size[4];
   pan_pack(wg_size, COMPUTE_SIZE_WORKGROUP, cfg) {
      cfg.workgroup_size_x = info->block[0];
      cfg.workgroup_size_y = info->block[1];
      cfg.workgroup_size_z = info->block[2];

      /* Workgroups may be merged if the shader does not use barriers
       * or shared memory. This condition is checked against the
       * static shared_size at compile-time. We need to check the
       * variable shared size at launch_grid time, because the
       * compiler doesn't know about that.
       */
      cfg.allow_merging_workgroups = cs->info.cs.allow_merging_workgroups &&
                                     (info->variable_shared_mem == 0);
   }

   cs_move32_to(b, cs_reg32(b, 33), wg_size[0]);

   /* Offset */
   for (unsigned i = 0; i < 3; ++i)
      cs_move32_to(b, cs_reg32(b, 34 + i), 0);

   unsigned threads_per_wg = info->block[0] * info->block[1] * info->block[2];
   unsigned max_thread_cnt = panfrost_compute_max_thread_count(
      &dev->kmod.props, cs->info.work_reg_count);

   if (info->indirect) {
      /* Load size in workgroups per dimension from memory */
      struct cs_index address = cs_reg64(b, 64);
      cs_move64_to(
         b, address,
         pan_resource(info->indirect)->image.data.base + info->indirect_offset);

      struct cs_index grid_xyz = cs_reg_tuple(b, 37, 3);
      cs_load_to(b, grid_xyz, address, BITFIELD_MASK(3), 0);

      /* Wait for the load */
      cs_wait_slot(b, 0, false);

      /* Copy to FAU */
      for (unsigned i = 0; i < 3; ++i) {
         if (batch->num_wg_sysval[i]) {
            cs_move64_to(b, address, batch->num_wg_sysval[i]);
            cs_store(b, cs_extract32(b, grid_xyz, i), address, BITFIELD_MASK(1),
                     0);
         }
      }

      /* Wait for the stores */
      cs_wait_slot(b, 0, false);

      cs_run_compute_indirect(b, DIV_ROUND_UP(max_thread_cnt, threads_per_wg),
                              false, cs_shader_res_sel(0, 0, 0, 0));
   } else {
      /* Set size in workgroups per dimension immediately */
      for (unsigned i = 0; i < 3; ++i)
         cs_move32_to(b, cs_reg32(b, 37 + i), info->grid[i]);

      /* Pick the task_axis and task_increment to maximize thread utilization. */
      unsigned task_axis = MALI_TASK_AXIS_X;
      unsigned threads_per_task = threads_per_wg;
      unsigned task_increment = 0;

      for (unsigned i = 0; i < 3; i++) {
         if (threads_per_task * info->grid[i] >= max_thread_cnt) {
            /* We reached out thread limit, stop at the current axis and
             * calculate the increment so it doesn't exceed the per-core
             * thread capacity.
             */
            task_increment = max_thread_cnt / threads_per_task;
            break;
         } else if (task_axis == MALI_TASK_AXIS_Z) {
            /* We reached the Z axis, and there's still room to stuff more
             * threads. Pick the current axis grid size as our increment
             * as there's no point using something bigger.
             */
            task_increment = info->grid[i];
            break;
         }

         threads_per_task *= info->grid[i];
         task_axis++;
      }

      assert(task_axis <= MALI_TASK_AXIS_Z);
      assert(task_increment > 0);
      cs_run_compute(b, task_increment, task_axis, false,
                     cs_shader_res_sel(0, 0, 0, 0));
   }
}

void
GENX(csf_launch_xfb)(struct panfrost_batch *batch,
                     const struct pipe_draw_info *info, unsigned count)
{
   struct cs_builder *b = batch->csf.cs.builder;

   cs_move64_to(b, cs_reg64(b, 24), batch->tls.gpu);

   /* TODO: Indexing. Also, attribute_offset is a legacy feature.. */
   cs_move32_to(b, cs_reg32(b, 32), batch->ctx->offset_start);

   /* Compute workgroup size */
   uint32_t wg_size[4];
   pan_pack(wg_size, COMPUTE_SIZE_WORKGROUP, cfg) {
      cfg.workgroup_size_x = 1;
      cfg.workgroup_size_y = 1;
      cfg.workgroup_size_z = 1;

      /* Transform feedback shaders do not use barriers or
       * shared memory, so we may merge workgroups.
       */
      cfg.allow_merging_workgroups = true;
   }
   cs_move32_to(b, cs_reg32(b, 33), wg_size[0]);

   /* Offset */
   for (unsigned i = 0; i < 3; ++i)
      cs_move32_to(b, cs_reg32(b, 34 + i), 0);

   cs_move32_to(b, cs_reg32(b, 37), count);
   cs_move32_to(b, cs_reg32(b, 38), info->instance_count);
   cs_move32_to(b, cs_reg32(b, 39), 1);

   csf_emit_shader_regs(batch, PIPE_SHADER_VERTEX,
                        batch->rsd[PIPE_SHADER_VERTEX]);
   /* force a barrier to avoid read/write sync issues with buffers */
   cs_wait_slot(b, 2, false);

   /* XXX: Choose correctly */
   cs_run_compute(b, 1, MALI_TASK_AXIS_Z, false, cs_shader_res_sel(0, 0, 0, 0));
}

static mali_ptr
csf_get_tiler_desc(struct panfrost_batch *batch)
{
   struct panfrost_context *ctx = batch->ctx;
   struct panfrost_device *dev = pan_device(ctx->base.screen);

   if (batch->tiler_ctx.valhall.desc)
      return batch->tiler_ctx.valhall.desc;

   struct panfrost_ptr t =
      pan_pool_alloc_desc(&batch->pool.base, TILER_CONTEXT);
   pan_pack(t.cpu, TILER_CONTEXT, tiler) {
      unsigned max_levels = dev->tiler_features.max_levels;
      assert(max_levels >= 2);

      /* TODO: Select hierarchy mask more effectively */
      tiler.hierarchy_mask = (max_levels >= 8) ? 0xFF : 0x28;

      /* For large framebuffers, disable the smallest bin size to
       * avoid pathological tiler memory usage. Required to avoid OOM
       * on dEQP-GLES31.functional.fbo.no_attachments.maximums.all on
       * Mali-G57.
       */
      if (MAX2(batch->key.width, batch->key.height) >= 4096)
         tiler.hierarchy_mask &= ~1;

      tiler.fb_width = batch->key.width;
      tiler.fb_height = batch->key.height;
      tiler.heap = batch->ctx->csf.heap.desc_bo->ptr.gpu;
      tiler.sample_pattern =
         pan_sample_pattern(util_framebuffer_get_num_samples(&batch->key));
      tiler.first_provoking_vertex =
         pan_tristate_get(batch->first_provoking_vertex);
      tiler.geometry_buffer = ctx->csf.tmp_geom_bo->ptr.gpu;
      tiler.geometry_buffer_size = ctx->csf.tmp_geom_bo->kmod_bo->size;
   }

   batch->tiler_ctx.valhall.desc = t.gpu;
   return batch->tiler_ctx.valhall.desc;
}

static uint32_t
csf_emit_draw_state(struct panfrost_batch *batch,
                    const struct pipe_draw_info *info, unsigned drawid_offset)
{
   struct panfrost_context *ctx = batch->ctx;
   struct panfrost_compiled_shader *vs = ctx->prog[PIPE_SHADER_VERTEX];
   struct panfrost_compiled_shader *fs = ctx->prog[PIPE_SHADER_FRAGMENT];

   bool idvs = vs->info.vs.idvs;
   bool fs_required = panfrost_fs_required(
      fs, ctx->blend, &ctx->pipe_framebuffer, ctx->depth_stencil);
   bool secondary_shader = vs->info.vs.secondary_enable && fs_required;

   assert(idvs && "IDVS required for CSF");

   struct cs_builder *b = batch->csf.cs.builder;

   if (batch->draw_count == 0)
      cs_vt_start(batch->csf.cs.builder, cs_now());

   csf_emit_shader_regs(batch, PIPE_SHADER_VERTEX,
                        panfrost_get_position_shader(batch, info));

   if (fs_required) {
      csf_emit_shader_regs(batch, PIPE_SHADER_FRAGMENT,
                           batch->rsd[PIPE_SHADER_FRAGMENT]);
   } else {
      cs_move64_to(b, cs_reg64(b, 4), 0);
      cs_move64_to(b, cs_reg64(b, 12), 0);
      cs_move64_to(b, cs_reg64(b, 20), 0);
   }

   if (secondary_shader) {
      cs_move64_to(b, cs_reg64(b, 18), panfrost_get_varying_shader(batch));
   }

   cs_move64_to(b, cs_reg64(b, 24), batch->tls.gpu);
   cs_move64_to(b, cs_reg64(b, 30), batch->tls.gpu);
   cs_move32_to(b, cs_reg32(b, 32), 0);
   cs_move32_to(b, cs_reg32(b, 37), 0);
   cs_move32_to(b, cs_reg32(b, 38), 0);

   cs_move64_to(b, cs_reg64(b, 40), csf_get_tiler_desc(batch));

   STATIC_ASSERT(sizeof(batch->scissor) == pan_size(SCISSOR));
   STATIC_ASSERT(sizeof(uint64_t) == pan_size(SCISSOR));
   uint64_t *sbd = (uint64_t *)&batch->scissor[0];
   cs_move64_to(b, cs_reg64(b, 42), *sbd);

   cs_move32_to(b, cs_reg32(b, 44), fui(batch->minimum_z));
   cs_move32_to(b, cs_reg32(b, 45), fui(batch->maximum_z));

   if (ctx->occlusion_query && ctx->active_queries) {
      struct panfrost_resource *rsrc = pan_resource(ctx->occlusion_query->rsrc);
      cs_move64_to(b, cs_reg64(b, 46), rsrc->image.data.base);
      panfrost_batch_write_rsrc(ctx->batch, rsrc, PIPE_SHADER_FRAGMENT);
   }

   cs_move32_to(b, cs_reg32(b, 48), panfrost_vertex_attribute_stride(vs, fs));
   cs_move64_to(b, cs_reg64(b, 50),
                batch->blend | MAX2(batch->key.nr_cbufs, 1));
   cs_move64_to(b, cs_reg64(b, 52), batch->depth_stencil);

   if (info->index_size)
      cs_move64_to(b, cs_reg64(b, 54), batch->indices);

   struct pipe_rasterizer_state *rast = &ctx->rasterizer->base;

   uint32_t primitive_flags = 0;
   pan_pack(&primitive_flags, PRIMITIVE_FLAGS, cfg) {
      if (panfrost_writes_point_size(ctx))
         cfg.point_size_array_format = MALI_POINT_SIZE_ARRAY_FORMAT_FP16;

      cfg.allow_rotating_primitives = allow_rotating_primitives(fs, info);

      cfg.low_depth_cull = rast->depth_clip_near;
      cfg.high_depth_cull = rast->depth_clip_far;

      /* Non-fixed restart indices should have been lowered */
      assert(!cfg.primitive_restart || panfrost_is_implicit_prim_restart(info));
      cfg.primitive_restart = info->primitive_restart;

      cfg.position_fifo_format = panfrost_writes_point_size(ctx)
                                    ? MALI_FIFO_FORMAT_EXTENDED
                                    : MALI_FIFO_FORMAT_BASIC;
   }

   cs_move32_to(b, cs_reg32(b, 56), primitive_flags);

   uint32_t dcd_flags0 = 0, dcd_flags1 = 0;
   pan_pack(&dcd_flags0, DCD_FLAGS_0, cfg) {
      enum mesa_prim reduced_mode = u_reduced_prim(info->mode);
      bool polygon = reduced_mode == MESA_PRIM_TRIANGLES;
      bool lines = reduced_mode == MESA_PRIM_LINES;

      /*
       * From the Gallium documentation,
       * pipe_rasterizer_state::cull_face "indicates which faces of
       * polygons to cull". Points and lines are not considered
       * polygons and should be drawn even if all faces are culled.
       * The hardware does not take primitive type into account when
       * culling, so we need to do that check ourselves.
       */
      cfg.cull_front_face = polygon && (rast->cull_face & PIPE_FACE_FRONT);
      cfg.cull_back_face = polygon && (rast->cull_face & PIPE_FACE_BACK);
      cfg.front_face_ccw = rast->front_ccw;

      cfg.multisample_enable = rast->multisample;

      /* Use per-sample shading if required by API Also use it when a
       * blend shader is used with multisampling, as this is handled
       * by a single ST_TILE in the blend shader with the current
       * sample ID, requiring per-sample shading.
       */
      cfg.evaluate_per_sample =
         (rast->multisample &&
          ((ctx->min_samples > 1) || ctx->valhall_has_blend_shader));

      cfg.single_sampled_lines = !rast->multisample;

      if (lines && rast->line_smooth) {
         cfg.multisample_enable = true;
         cfg.single_sampled_lines = false;
      }

      bool has_oq = ctx->occlusion_query && ctx->active_queries;
      if (has_oq) {
         if (ctx->occlusion_query->type == PIPE_QUERY_OCCLUSION_COUNTER)
            cfg.occlusion_query = MALI_OCCLUSION_MODE_COUNTER;
         else
            cfg.occlusion_query = MALI_OCCLUSION_MODE_PREDICATE;
      }

      if (fs_required) {
         struct pan_earlyzs_state earlyzs = pan_earlyzs_get(
            fs->earlyzs, ctx->depth_stencil->writes_zs || has_oq,
            ctx->blend->base.alpha_to_coverage,
            ctx->depth_stencil->zs_always_passes);

         cfg.pixel_kill_operation = earlyzs.kill;
         cfg.zs_update_operation = earlyzs.update;

         cfg.allow_forward_pixel_to_kill =
            pan_allow_forward_pixel_to_kill(ctx, fs);
         cfg.allow_forward_pixel_to_be_killed = !fs->info.writes_global;

         cfg.overdraw_alpha0 = panfrost_overdraw_alpha(ctx, 0);
         cfg.overdraw_alpha1 = panfrost_overdraw_alpha(ctx, 1);

         /* Also use per-sample shading if required by the shader
          */
         cfg.evaluate_per_sample |= fs->info.fs.sample_shading;

         /* Unlike Bifrost, alpha-to-coverage must be included in
          * this identically-named flag. Confusing, isn't it?
          */
         cfg.shader_modifies_coverage = fs->info.fs.writes_coverage ||
                                        fs->info.fs.can_discard ||
                                        ctx->blend->base.alpha_to_coverage;

         cfg.alpha_to_coverage = ctx->blend->base.alpha_to_coverage;
      } else {
         /* These operations need to be FORCE to benefit from the
          * depth-only pass optimizations.
          */
         cfg.pixel_kill_operation = MALI_PIXEL_KILL_FORCE_EARLY;
         cfg.zs_update_operation = MALI_PIXEL_KILL_FORCE_EARLY;

         /* No shader and no blend => no shader or blend
          * reasons to disable FPK. The only FPK-related state
          * not covered is alpha-to-coverage which we don't set
          * without blend.
          */
         cfg.allow_forward_pixel_to_kill = true;

         /* No shader => no shader side effects */
         cfg.allow_forward_pixel_to_be_killed = true;

         /* Alpha isn't written so these are vacuous */
         cfg.overdraw_alpha0 = true;
         cfg.overdraw_alpha1 = true;
      }
   }

   pan_pack(&dcd_flags1, DCD_FLAGS_1, cfg) {
      cfg.sample_mask = rast->multisample ? ctx->sample_mask : 0xFFFF;

      if (fs_required) {
         /* See JM Valhall equivalent code */
         cfg.render_target_mask =
            (fs->info.outputs_written >> FRAG_RESULT_DATA0) & ctx->fb_rt_mask;
      }
   }

   cs_move32_to(b, cs_reg32(b, 57), dcd_flags0);
   cs_move32_to(b, cs_reg32(b, 58), dcd_flags1);

   uint64_t primsize = 0;
   panfrost_emit_primitive_size(ctx, info->mode == MESA_PRIM_POINTS, 0,
                                &primsize);
   cs_move64_to(b, cs_reg64(b, 60), primsize);

   uint32_t flags_override;
   /* Pack with nodefaults so only explicitly set override fields affect the
    * previously set register values */
   pan_pack_nodefaults(&flags_override, PRIMITIVE_FLAGS, cfg) {
      cfg.draw_mode = pan_draw_mode(info->mode);
      cfg.index_type = panfrost_translate_index_size(info->index_size);
      cfg.secondary_shader = secondary_shader;
   };

   return flags_override;
}

static struct cs_index
csf_emit_draw_id_register(struct panfrost_batch *batch, unsigned offset)
{
   struct cs_builder *b = batch->csf.cs.builder;
   struct panfrost_context *ctx = batch->ctx;
   struct panfrost_uncompiled_shader *vs = ctx->uncompiled[PIPE_SHADER_VERTEX];

   if (!BITSET_TEST(vs->nir->info.system_values_read, SYSTEM_VALUE_DRAW_ID))
      return cs_undef();

   struct cs_index drawid = cs_reg32(b, 67);
   cs_move32_to(b, drawid, offset);

   return drawid;
}

void
GENX(csf_launch_draw)(struct panfrost_batch *batch,
                      const struct pipe_draw_info *info, unsigned drawid_offset,
                      const struct pipe_draw_start_count_bias *draw,
                      unsigned vertex_count)
{
   struct cs_builder *b = batch->csf.cs.builder;

   uint32_t flags_override = csf_emit_draw_state(batch, info, drawid_offset);
   struct cs_index drawid = csf_emit_draw_id_register(batch, drawid_offset);

   cs_move32_to(b, cs_reg32(b, 33), draw->count);
   cs_move32_to(b, cs_reg32(b, 34), info->instance_count);
   cs_move32_to(b, cs_reg32(b, 35), 0);

   /* Base vertex offset on Valhall is used for both indexed and
    * non-indexed draws, in a simple way for either. Handle both cases.
    */
   if (info->index_size) {
      cs_move32_to(b, cs_reg32(b, 36), draw->index_bias);
      cs_move32_to(b, cs_reg32(b, 39), info->index_size * draw->count);
   } else {
      cs_move32_to(b, cs_reg32(b, 36), draw->start);
      cs_move32_to(b, cs_reg32(b, 39), 0);
   }

   cs_run_idvs(b, flags_override, false, true, cs_shader_res_sel(0, 0, 1, 0),
               cs_shader_res_sel(2, 2, 2, 0), drawid);
}

void
GENX(csf_launch_draw_indirect)(struct panfrost_batch *batch,
                               const struct pipe_draw_info *info,
                               unsigned drawid_offset,
                               const struct pipe_draw_indirect_info *indirect)
{
   struct cs_builder *b = batch->csf.cs.builder;

   uint32_t flags_override = csf_emit_draw_state(batch, info, drawid_offset);
   struct cs_index drawid = csf_emit_draw_id_register(batch, drawid_offset);

   struct cs_index address = cs_reg64(b, 64);
   struct cs_index counter = cs_reg32(b, 66);
   cs_move64_to(
      b, address,
      pan_resource(indirect->buffer)->image.data.base + indirect->offset);
   cs_move32_to(b, counter, indirect->draw_count);

   cs_while(b, MALI_CS_CONDITION_GREATER, counter) {
      if (info->index_size) {
         /* loads vertex count, instance count, index offset, vertex offset */
         cs_load_to(b, cs_reg_tuple(b, 33, 4), address, BITFIELD_MASK(4), 0);
         cs_move32_to(b, cs_reg32(b, 39), info->index.resource->width0);
      } else {
         /* vertex count, instance count */
         cs_load_to(b, cs_reg_tuple(b, 33, 2), address, BITFIELD_MASK(2), 0);
         cs_move32_to(b, cs_reg32(b, 35), 0);
         cs_load_to(b, cs_reg_tuple(b, 36, 1), address, BITFIELD_MASK(1),
                    2 * sizeof(uint32_t)); // instance offset
         cs_move32_to(b, cs_reg32(b, 37), 0);
         cs_move32_to(b, cs_reg32(b, 39), 0);
      }

      cs_wait_slot(b, 0, false);
      cs_run_idvs(b, flags_override, false, true, cs_shader_res_sel(0, 0, 1, 0),
                  cs_shader_res_sel(2, 2, 2, 0), drawid);

      cs_add64(b, address, address, indirect->stride);
      cs_add32(b, counter, counter, (unsigned int)-1);
      if (drawid.type != CS_INDEX_UNDEF)
         cs_add32(b, drawid, drawid, 1);
   }
}

#define POSITION_FIFO_SIZE (64 * 1024)

int
GENX(csf_init_context)(struct panfrost_context *ctx)
{
   struct panfrost_device *dev = pan_device(ctx->base.screen);
   struct drm_panthor_queue_create qc[] = {{
      .priority = 1,
      .ringbuf_size = 64 * 1024,
   }};

   struct drm_panthor_group_create gc = {
      .compute_core_mask = dev->kmod.props.shader_present,
      .fragment_core_mask = dev->kmod.props.shader_present,
      .tiler_core_mask = 1,
      .max_compute_cores = util_bitcount64(dev->kmod.props.shader_present),
      .max_fragment_cores = util_bitcount64(dev->kmod.props.shader_present),
      .max_tiler_cores = 1,
      .priority = PANTHOR_GROUP_PRIORITY_MEDIUM,
      .queues = DRM_PANTHOR_OBJ_ARRAY(ARRAY_SIZE(qc), qc),
      .vm_id = pan_kmod_vm_handle(dev->kmod.vm),
   };

   int ret =
      drmIoctl(panfrost_device_fd(dev), DRM_IOCTL_PANTHOR_GROUP_CREATE, &gc);

   if (ret)
      goto err_group_create;

   ctx->csf.group_handle = gc.group_handle;

   struct drm_panthor_group_destroy gd = {
      .group_handle = ctx->csf.group_handle,
   };

   /* Get tiler heap */
   struct drm_panthor_tiler_heap_create thc = {
      .vm_id = pan_kmod_vm_handle(dev->kmod.vm),
      .chunk_size = pan_screen(ctx->base.screen)->csf_tiler_heap.chunk_size,
      .initial_chunk_count =
         pan_screen(ctx->base.screen)->csf_tiler_heap.initial_chunks,
      .max_chunks = pan_screen(ctx->base.screen)->csf_tiler_heap.max_chunks,
      .target_in_flight = 65535,
   };
   ret = drmIoctl(panfrost_device_fd(dev), DRM_IOCTL_PANTHOR_TILER_HEAP_CREATE,
                  &thc);

   if (ret)
      goto err_tiler_heap;

   ctx->csf.heap.handle = thc.handle;

   struct drm_panthor_tiler_heap_destroy thd = {
      .handle = ctx->csf.heap.handle,
   };

   ctx->csf.heap.desc_bo =
      panfrost_bo_create(dev, pan_size(TILER_HEAP), 0, "Tiler Heap");

   if (ctx->csf.heap.desc_bo == NULL)
      goto err_tiler_heap_desc_bo;

   pan_pack(ctx->csf.heap.desc_bo->ptr.cpu, TILER_HEAP, heap) {
      heap.size = pan_screen(ctx->base.screen)->csf_tiler_heap.chunk_size;
      heap.base = thc.first_heap_chunk_gpu_va;
      heap.bottom = heap.base + 64;
      heap.top = heap.base + heap.size;
   }

   ctx->csf.tmp_geom_bo = panfrost_bo_create(
      dev, POSITION_FIFO_SIZE, PAN_BO_INVISIBLE, "Temporary Geometry buffer");

   if (ctx->csf.tmp_geom_bo == NULL)
      goto err_tiler_heap_tmp_geom_bo;

   /* Setup the tiler heap */
   struct panfrost_bo *cs_bo =
      panfrost_bo_create(dev, 4096, 0, "Temporary CS buffer");

   if (cs_bo == NULL)
      goto err_tiler_heap_cs_bo;

   struct cs_buffer init_buffer = {
      .cpu = cs_bo->ptr.cpu,
      .gpu = cs_bo->ptr.gpu,
      .capacity = panfrost_bo_size(cs_bo) / sizeof(uint64_t),
   };
   const struct cs_builder_conf bconf = {
      .nr_registers = 96,
      .nr_kernel_registers = 4,
   };
   struct cs_builder b;
   cs_builder_init(&b, &bconf, init_buffer);
   struct cs_index heap = cs_reg64(&b, 72);
   cs_move64_to(&b, heap, thc.tiler_heap_ctx_gpu_va);
   cs_heap_set(&b, heap);

   struct drm_panthor_queue_submit qsubmit;
   struct drm_panthor_group_submit gsubmit;
   struct drm_panthor_sync_op sync = {
      .flags =
         DRM_PANTHOR_SYNC_OP_SIGNAL | DRM_PANTHOR_SYNC_OP_HANDLE_TYPE_SYNCOBJ,
      .handle = ctx->syncobj,
   };

   assert(cs_is_valid(&b));
   cs_finish(&b);

   uint64_t cs_start = cs_root_chunk_gpu_addr(&b);
   uint32_t cs_size = cs_root_chunk_size(&b);

   csf_prepare_qsubmit(ctx, &qsubmit, 0, cs_start, cs_size, &sync, 1);
   csf_prepare_gsubmit(ctx, &gsubmit, &qsubmit, 1);
   ret = csf_submit_gsubmit(ctx, &gsubmit);

   if (ret)
      goto err_g_submit;

   /* Wait before freeing the buffer. */
   ret = drmSyncobjWait(panfrost_device_fd(dev), &ctx->syncobj, 1, INT64_MAX, 0,
                        NULL);
   assert(!ret);

   panfrost_bo_unreference(cs_bo);

   ctx->csf.is_init = true;
   return 0;
err_g_submit:
   panfrost_bo_unreference(cs_bo);
err_tiler_heap_cs_bo:
   panfrost_bo_unreference(ctx->csf.tmp_geom_bo);
err_tiler_heap_tmp_geom_bo:
   panfrost_bo_unreference(ctx->csf.heap.desc_bo);
err_tiler_heap_desc_bo:
   drmIoctl(panfrost_device_fd(dev), DRM_IOCTL_PANTHOR_TILER_HEAP_DESTROY,
            &thd);
err_tiler_heap:
   drmIoctl(panfrost_device_fd(dev), DRM_IOCTL_PANTHOR_GROUP_DESTROY, &gd);
err_group_create:
   return -1;
}

void
GENX(csf_cleanup_context)(struct panfrost_context *ctx)
{
   if (!ctx->csf.is_init)
      return;

   struct panfrost_device *dev = pan_device(ctx->base.screen);
   struct drm_panthor_tiler_heap_destroy thd = {
      .handle = ctx->csf.heap.handle,
   };
   int ret;

   /* Make sure all jobs are done before destroying the heap. */
   ret = drmSyncobjWait(panfrost_device_fd(dev), &ctx->syncobj, 1, INT64_MAX, 0,
                        NULL);
   assert(!ret);

   ret = drmIoctl(panfrost_device_fd(dev), DRM_IOCTL_PANTHOR_TILER_HEAP_DESTROY,
                  &thd);
   assert(!ret);

   struct drm_panthor_group_destroy gd = {
      .group_handle = ctx->csf.group_handle,
   };

   ret =
      drmIoctl(panfrost_device_fd(dev), DRM_IOCTL_PANTHOR_GROUP_DESTROY, &gd);
   assert(!ret);

   panfrost_bo_unreference(ctx->csf.heap.desc_bo);
   ctx->csf.is_init = false;
}

void
GENX(csf_emit_write_timestamp)(struct panfrost_batch *batch,
                               struct panfrost_resource *dst, unsigned offset)
{
   struct cs_builder *b = batch->csf.cs.builder;

   struct cs_index address = cs_reg64(b, 40);
   cs_move64_to(b, address,
                dst->image.data.base + dst->image.data.offset + offset);
   cs_store_state(b, address, 0, MALI_CS_STATE_TIMESTAMP, cs_now());

   panfrost_batch_write_rsrc(batch, dst, PIPE_SHADER_VERTEX);
}