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

/*
 * Copyright (C) 2018 Alyssa Rosenzweig
 * Copyright (C) 2020 Collabora Ltd.
 * Copyright © 2017 Intel Corporation
 *
 * 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/panfrost_drm.h"

#include "pan_blitter.h"
#include "pan_cmdstream.h"
#include "pan_context.h"
#include "pan_indirect_dispatch.h"
#include "pan_jm.h"
#include "pan_job.h"

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

void
GENX(jm_init_batch)(struct panfrost_batch *batch)
{
   /* Reserve the framebuffer and local storage descriptors */
   batch->framebuffer =
#if PAN_ARCH == 4
      pan_pool_alloc_desc(&batch->pool.base, FRAMEBUFFER);
#else
      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));
#endif

#if PAN_ARCH >= 6
   batch->tls = pan_pool_alloc_desc(&batch->pool.base, LOCAL_STORAGE);
#else
   /* On Midgard, the TLS is embedded in the FB descriptor */
   batch->tls = batch->framebuffer;

#if PAN_ARCH == 5
   struct mali_framebuffer_pointer_packed ptr;

   pan_pack(ptr.opaque, FRAMEBUFFER_POINTER, cfg) {
      cfg.pointer = batch->framebuffer.gpu;
      cfg.render_target_count = 1; /* a necessary lie */
   }

   batch->tls.gpu = ptr.opaque[0];
#endif
#endif
}

static int
jm_submit_jc(struct panfrost_batch *batch, mali_ptr first_job_desc,
             uint32_t reqs, uint32_t out_sync)
{
   struct panfrost_context *ctx = batch->ctx;
   struct pipe_context *gallium = (struct pipe_context *)ctx;
   struct panfrost_device *dev = pan_device(gallium->screen);
   struct drm_panfrost_submit submit = {
      0,
   };
   uint32_t in_syncs[1];
   uint32_t *bo_handles;
   int ret;

   /* If we trace, we always need a syncobj, so make one of our own if we
    * weren't given one to use. Remember that we did so, so we can free it
    * after we're done but preventing double-frees if we were given a
    * syncobj */

   if (!out_sync && dev->debug & (PAN_DBG_TRACE | PAN_DBG_SYNC))
      out_sync = ctx->syncobj;

   submit.out_sync = out_sync;
   submit.jc = first_job_desc;
   submit.requirements = reqs;

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

      in_syncs[submit.in_sync_count++] = ctx->in_sync_obj;
      close(ctx->in_sync_fd);
      ctx->in_sync_fd = -1;
   }

   if (submit.in_sync_count)
      submit.in_syncs = (uintptr_t)in_syncs;

   bo_handles = calloc(panfrost_pool_num_bos(&batch->pool) +
                          panfrost_pool_num_bos(&batch->invisible_pool) +
                          batch->num_bos + 2,
                       sizeof(*bo_handles));
   assert(bo_handles);

   pan_bo_access *flags = util_dynarray_begin(&batch->bos);
   unsigned end_bo = util_dynarray_num_elements(&batch->bos, pan_bo_access);

   for (int i = 0; i < end_bo; ++i) {
      if (!flags[i])
         continue;

      assert(submit.bo_handle_count < batch->num_bos);
      bo_handles[submit.bo_handle_count++] = i;

      /* 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);

      bo->gpu_access |= flags[i] & (PAN_BO_ACCESS_RW);
   }

   panfrost_pool_get_bo_handles(&batch->pool,
                                bo_handles + submit.bo_handle_count);
   submit.bo_handle_count += panfrost_pool_num_bos(&batch->pool);
   panfrost_pool_get_bo_handles(&batch->invisible_pool,
                                bo_handles + submit.bo_handle_count);
   submit.bo_handle_count += panfrost_pool_num_bos(&batch->invisible_pool);

   /* Add the tiler heap to the list of accessed BOs if the batch has at
    * least one tiler job. Tiler heap is written by tiler jobs and read
    * by fragment jobs (the polygon list is coming from this heap).
    */
   if (batch->jm.jobs.vtc_jc.first_tiler)
      bo_handles[submit.bo_handle_count++] =
         panfrost_bo_handle(dev->tiler_heap);

   /* Always used on Bifrost, occassionally used on Midgard */
   bo_handles[submit.bo_handle_count++] =
      panfrost_bo_handle(dev->sample_positions);

   submit.bo_handles = (u64)(uintptr_t)bo_handles;
   if (ctx->is_noop)
      ret = 0;
   else
      ret = drmIoctl(panfrost_device_fd(dev), DRM_IOCTL_PANFROST_SUBMIT, &submit);
   free(bo_handles);

   if (ret)
      return errno;

   /* Trace the job if we're doing that */
   if (dev->debug & (PAN_DBG_TRACE | PAN_DBG_SYNC)) {
      /* Wait so we can get errors reported back */
      ret = drmSyncobjWait(panfrost_device_fd(dev), &out_sync, 1, INT64_MAX,
                           0, NULL);
      if (ret)
         return errno;

      if (dev->debug & PAN_DBG_TRACE)
         pandecode_jc(dev->decode_ctx, submit.jc, panfrost_device_gpu_id(dev));

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

      /* Jobs won't be complete if blackhole rendering, that's ok */
      if (!ctx->is_noop && dev->debug & PAN_DBG_SYNC)
         pandecode_abort_on_fault(dev->decode_ctx, submit.jc, panfrost_device_gpu_id(dev));
   }

   return 0;
}

/* Submit both vertex/tiler and fragment jobs for a batch, possibly with an
 * outsync corresponding to the later of the two (since there will be an
 * implicit dep between them) */

int
GENX(jm_submit_batch)(struct panfrost_batch *batch)
{
   struct pipe_screen *pscreen = batch->ctx->base.screen;
   struct panfrost_device *dev = pan_device(pscreen);
   bool has_draws = batch->jm.jobs.vtc_jc.first_job;
   bool has_tiler = batch->jm.jobs.vtc_jc.first_tiler;
   bool has_frag = panfrost_has_fragment_job(batch);
   uint32_t out_sync = batch->ctx->syncobj;
   int ret = 0;

   unsigned reqs =
      batch->need_job_req_cycle_count ? PANFROST_JD_REQ_CYCLE_COUNT : 0;

   /* Take the submit lock to make sure no tiler jobs from other context
    * are inserted between our tiler and fragment jobs, failing to do that
    * might result in tiler heap corruption.
    */
   if (has_tiler)
      pthread_mutex_lock(&dev->submit_lock);

   if (has_draws) {
      ret = jm_submit_jc(batch, batch->jm.jobs.vtc_jc.first_job, reqs,
                         has_frag ? 0 : out_sync);

      if (ret)
         goto done;
   }

   if (has_frag) {
      ret = jm_submit_jc(batch, batch->jm.jobs.frag, reqs | PANFROST_JD_REQ_FS,
                         out_sync);
      if (ret)
         goto done;
   }

done:
   if (has_tiler)
      pthread_mutex_unlock(&dev->submit_lock);

   return ret;
}

void
GENX(jm_preload_fb)(struct panfrost_batch *batch, struct pan_fb_info *fb)
{
   struct panfrost_device *dev = pan_device(batch->ctx->base.screen);
   struct panfrost_ptr preload_jobs[2];

   unsigned preload_job_count = GENX(pan_preload_fb)(
      &dev->blitter, &batch->pool.base, fb, 0, batch->tls.gpu, preload_jobs);

   assert(PAN_ARCH < 6 || !preload_job_count);

   for (unsigned j = 0; j < preload_job_count; j++) {
      pan_jc_add_job(&batch->jm.jobs.vtc_jc, MALI_JOB_TYPE_TILER, false, false,
                     0, 0, &preload_jobs[j], true);
   }
}

void
GENX(jm_emit_fragment_job)(struct panfrost_batch *batch,
                           const struct pan_fb_info *pfb)
{
   struct panfrost_ptr transfer =
      pan_pool_alloc_desc(&batch->pool.base, FRAGMENT_JOB);

   GENX(pan_emit_fragment_job_payload)
   (pfb, batch->framebuffer.gpu, transfer.cpu);

   pan_section_pack(transfer.cpu, FRAGMENT_JOB, HEADER, header) {
      header.type = MALI_JOB_TYPE_FRAGMENT;
      header.index = 1;
   }

   batch->jm.jobs.frag = transfer.gpu;
}

#if PAN_ARCH == 9
static void
jm_emit_shader_env(struct panfrost_batch *batch,
                   struct MALI_SHADER_ENVIRONMENT *cfg,
                   enum pipe_shader_type stage, mali_ptr shader_ptr)
{
   cfg->resources = panfrost_emit_resources(batch, stage);
   cfg->thread_storage = batch->tls.gpu;
   cfg->shader = shader_ptr;

   /* Each entry of FAU is 64-bits */
   cfg->fau = batch->push_uniforms[stage];
   cfg->fau_count = DIV_ROUND_UP(batch->nr_push_uniforms[stage], 2);
}
#endif

void
GENX(jm_launch_grid)(struct panfrost_batch *batch,
                     const struct pipe_grid_info *info)
{
   struct panfrost_ptr t = pan_pool_alloc_desc(&batch->pool.base, COMPUTE_JOB);

   /* Invoke according to the grid info */

   unsigned num_wg[3] = {info->grid[0], info->grid[1], info->grid[2]};

   if (info->indirect)
      num_wg[0] = num_wg[1] = num_wg[2] = 1;

#if PAN_ARCH <= 7
   panfrost_pack_work_groups_compute(
      pan_section_ptr(t.cpu, COMPUTE_JOB, INVOCATION), num_wg[0], num_wg[1],
      num_wg[2], info->block[0], info->block[1], info->block[2], false,
      info->indirect != NULL);

   pan_section_pack(t.cpu, COMPUTE_JOB, PARAMETERS, cfg) {
      cfg.job_task_split = util_logbase2_ceil(info->block[0] + 1) +
                           util_logbase2_ceil(info->block[1] + 1) +
                           util_logbase2_ceil(info->block[2] + 1);
   }

   pan_section_pack(t.cpu, COMPUTE_JOB, DRAW, cfg) {
      cfg.state = batch->rsd[PIPE_SHADER_COMPUTE];
      cfg.attributes = batch->attribs[PIPE_SHADER_COMPUTE];
      cfg.attribute_buffers = batch->attrib_bufs[PIPE_SHADER_COMPUTE];
      cfg.thread_storage = batch->tls.gpu;
      cfg.uniform_buffers = batch->uniform_buffers[PIPE_SHADER_COMPUTE];
      cfg.push_uniforms = batch->push_uniforms[PIPE_SHADER_COMPUTE];
      cfg.textures = batch->textures[PIPE_SHADER_COMPUTE];
      cfg.samplers = batch->samplers[PIPE_SHADER_COMPUTE];
   }

#if PAN_ARCH == 4
   pan_section_pack(t.cpu, COMPUTE_JOB, COMPUTE_PADDING, cfg)
      ;
#endif
#else
   struct panfrost_context *ctx = batch->ctx;
   struct panfrost_compiled_shader *cs = ctx->prog[PIPE_SHADER_COMPUTE];

   pan_section_pack(t.cpu, COMPUTE_JOB, PAYLOAD, cfg) {
      cfg.workgroup_size_x = info->block[0];
      cfg.workgroup_size_y = info->block[1];
      cfg.workgroup_size_z = info->block[2];

      cfg.workgroup_count_x = num_wg[0];
      cfg.workgroup_count_y = num_wg[1];
      cfg.workgroup_count_z = num_wg[2];

      jm_emit_shader_env(batch, &cfg.compute, PIPE_SHADER_COMPUTE,
                         batch->rsd[PIPE_SHADER_COMPUTE]);

      /* 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);

      cfg.task_increment = 1;
      cfg.task_axis = MALI_TASK_AXIS_Z;
   }
#endif

   unsigned indirect_dep = 0;
#if PAN_GPU_SUPPORTS_DISPATCH_INDIRECT
   if (info->indirect) {
      struct panfrost_device *dev = pan_device(batch->ctx->base.screen);
      struct pan_indirect_dispatch_info indirect = {
         .job = t.gpu,
         .indirect_dim = pan_resource(info->indirect)->image.data.base +
                         info->indirect_offset,
         .num_wg_sysval =
            {
               batch->num_wg_sysval[0],
               batch->num_wg_sysval[1],
               batch->num_wg_sysval[2],
            },
      };

      indirect_dep = GENX(pan_indirect_dispatch_emit)(
         &dev->indirect_dispatch, &batch->pool.base, &batch->jm.jobs.vtc_jc,
         &indirect);
   }
#endif

   pan_jc_add_job(&batch->jm.jobs.vtc_jc, MALI_JOB_TYPE_COMPUTE, true, false,
                  indirect_dep, 0, &t, false);
}

#if PAN_ARCH >= 6
static mali_ptr
jm_emit_tiler_desc(struct panfrost_batch *batch)
{
   struct panfrost_device *dev = pan_device(batch->ctx->base.screen);
   mali_ptr tiler_desc = PAN_ARCH >= 9 ? batch->tiler_ctx.bifrost.desc
                                       : batch->tiler_ctx.valhall.desc;

   if (tiler_desc)
      return tiler_desc;

   struct panfrost_ptr t = pan_pool_alloc_desc(&batch->pool.base, TILER_HEAP);

   pan_pack(t.cpu, TILER_HEAP, heap) {
      heap.size = panfrost_bo_size(dev->tiler_heap);
      heap.base = dev->tiler_heap->ptr.gpu;
      heap.bottom = dev->tiler_heap->ptr.gpu;
      heap.top = dev->tiler_heap->ptr.gpu + panfrost_bo_size(dev->tiler_heap);
   }

   mali_ptr heap = t.gpu;
   unsigned max_levels = dev->tiler_features.max_levels;
   assert(max_levels >= 2);

   t = pan_pool_alloc_desc(&batch->pool.base, TILER_CONTEXT);
   pan_pack(t.cpu, TILER_CONTEXT, tiler) {
      /* 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 = heap;
      tiler.sample_pattern =
         pan_sample_pattern(util_framebuffer_get_num_samples(&batch->key));
#if PAN_ARCH >= 9
      tiler.first_provoking_vertex =
         pan_tristate_get(batch->first_provoking_vertex);
#endif
   }

   if (PAN_ARCH >= 9)
      batch->tiler_ctx.valhall.desc = t.gpu;
   else
      batch->tiler_ctx.bifrost.desc = t.gpu;

   return t.gpu;
}
#endif

#if PAN_ARCH <= 7
static inline void
jm_emit_draw_descs(struct panfrost_batch *batch, struct MALI_DRAW *d,
                   enum pipe_shader_type st)
{
   d->offset_start = batch->ctx->offset_start;
   d->instance_size =
      batch->ctx->instance_count > 1 ? batch->ctx->padded_count : 1;

   d->uniform_buffers = batch->uniform_buffers[st];
   d->push_uniforms = batch->push_uniforms[st];
   d->textures = batch->textures[st];
   d->samplers = batch->samplers[st];
}

static void
jm_emit_vertex_draw(struct panfrost_batch *batch, void *section)
{
   pan_pack(section, DRAW, cfg) {
      cfg.state = batch->rsd[PIPE_SHADER_VERTEX];
      cfg.attributes = batch->attribs[PIPE_SHADER_VERTEX];
      cfg.attribute_buffers = batch->attrib_bufs[PIPE_SHADER_VERTEX];
      cfg.varyings = batch->varyings.vs;
      cfg.varying_buffers = cfg.varyings ? batch->varyings.bufs : 0;
      cfg.thread_storage = batch->tls.gpu;
      jm_emit_draw_descs(batch, &cfg, PIPE_SHADER_VERTEX);
   }
}

static void
jm_emit_vertex_job(struct panfrost_batch *batch,
                   const struct pipe_draw_info *info, void *invocation_template,
                   void *job)
{
   void *section = pan_section_ptr(job, COMPUTE_JOB, INVOCATION);
   memcpy(section, invocation_template, pan_size(INVOCATION));

   pan_section_pack(job, COMPUTE_JOB, PARAMETERS, cfg) {
      cfg.job_task_split = 5;
   }

   section = pan_section_ptr(job, COMPUTE_JOB, DRAW);
   jm_emit_vertex_draw(batch, section);

#if PAN_ARCH == 4
   pan_section_pack(job, COMPUTE_JOB, COMPUTE_PADDING, cfg)
      ;
#endif
}
#endif /* PAN_ARCH <= 7 */

static void
jm_emit_tiler_draw(void *out, struct panfrost_batch *batch, bool fs_required,
                   enum mesa_prim prim)
{
   struct panfrost_context *ctx = batch->ctx;
   struct pipe_rasterizer_state *rast = &ctx->rasterizer->base;
   bool polygon = (prim == MESA_PRIM_TRIANGLES);

   pan_pack(out, DRAW, cfg) {
      /*
       * 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;

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

         struct panfrost_resource *rsrc =
            pan_resource(ctx->occlusion_query->rsrc);
         cfg.occlusion = rsrc->image.data.base;
         panfrost_batch_write_rsrc(ctx->batch, rsrc, PIPE_SHADER_FRAGMENT);
      }

#if PAN_ARCH >= 9
      struct panfrost_compiled_shader *fs = ctx->prog[PIPE_SHADER_FRAGMENT];

      cfg.multisample_enable = rast->multisample;
      cfg.sample_mask = rast->multisample ? ctx->sample_mask : 0xFFFF;

      /* 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;

      cfg.vertex_array.packet = true;

      cfg.minimum_z = batch->minimum_z;
      cfg.maximum_z = batch->maximum_z;

      cfg.depth_stencil = batch->depth_stencil;

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

      if (fs_required) {
         bool has_oq = ctx->occlusion_query && ctx->active_queries;

         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;

         /* Mask of render targets that may be written. A render
          * target may be written if the fragment shader writes
          * to it AND it actually exists. If the render target
          * doesn't actually exist, the blend descriptor will be
          * OFF so it may be omitted from the mask.
          *
          * Only set when there is a fragment shader, since
          * otherwise no colour updates are possible.
          */
         cfg.render_target_mask =
            (fs->info.outputs_written >> FRAG_RESULT_DATA0) & ctx->fb_rt_mask;

         /* 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;

         /* Blend descriptors are only accessed by a BLEND
          * instruction on Valhall. It follows that if the
          * fragment shader is omitted, we may also emit the
          * blend descriptors.
          */
         cfg.blend = batch->blend;
         cfg.blend_count = MAX2(batch->key.nr_cbufs, 1);
         cfg.alpha_to_coverage = ctx->blend->base.alpha_to_coverage;

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

         jm_emit_shader_env(batch, &cfg.shader, PIPE_SHADER_FRAGMENT,
                            batch->rsd[PIPE_SHADER_FRAGMENT]);
      } 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;
      }
#else
      cfg.position = batch->varyings.pos;
      cfg.state = batch->rsd[PIPE_SHADER_FRAGMENT];
      cfg.attributes = batch->attribs[PIPE_SHADER_FRAGMENT];
      cfg.attribute_buffers = batch->attrib_bufs[PIPE_SHADER_FRAGMENT];
      cfg.viewport = batch->viewport;
      cfg.varyings = batch->varyings.fs;
      cfg.varying_buffers = cfg.varyings ? batch->varyings.bufs : 0;
      cfg.thread_storage = batch->tls.gpu;

      /* For all primitives but lines DRAW.flat_shading_vertex must
       * be set to 0 and the provoking vertex is selected with the
       * PRIMITIVE.first_provoking_vertex field.
       */
      if (prim == MESA_PRIM_LINES) {
         /* The logic is inverted across arches. */
         cfg.flat_shading_vertex = rast->flatshade_first ^ (PAN_ARCH <= 5);
      }

      jm_emit_draw_descs(batch, &cfg, PIPE_SHADER_FRAGMENT);
#endif
   }
}

/* Packs a primitive descriptor, mostly common between Midgard/Bifrost tiler
 * jobs and Valhall IDVS jobs
 */
static void
jm_emit_primitive(struct panfrost_batch *batch,
                  const struct pipe_draw_info *info,
                  const struct pipe_draw_start_count_bias *draw,
                  bool secondary_shader, void *out)
{
   struct panfrost_context *ctx = batch->ctx;
   UNUSED struct pipe_rasterizer_state *rast = &ctx->rasterizer->base;

   pan_pack(out, PRIMITIVE, cfg) {
      cfg.draw_mode = pan_draw_mode(info->mode);
      if (panfrost_writes_point_size(ctx))
         cfg.point_size_array_format = MALI_POINT_SIZE_ARRAY_FORMAT_FP16;

#if PAN_ARCH <= 8
      /* For line primitives, PRIMITIVE.first_provoking_vertex must
       * be set to true and the provoking vertex is selected with
       * DRAW.flat_shading_vertex.
       */
      if (u_reduced_prim(info->mode) == MESA_PRIM_LINES)
         cfg.first_provoking_vertex = true;
      else
         cfg.first_provoking_vertex = rast->flatshade_first;

      if (panfrost_is_implicit_prim_restart(info)) {
         cfg.primitive_restart = MALI_PRIMITIVE_RESTART_IMPLICIT;
      } else if (info->primitive_restart) {
         cfg.primitive_restart = MALI_PRIMITIVE_RESTART_EXPLICIT;
         cfg.primitive_restart_index = info->restart_index;
      }

      cfg.job_task_split = 6;
#else
      struct panfrost_compiled_shader *fs = ctx->prog[PIPE_SHADER_FRAGMENT];

      cfg.allow_rotating_primitives = allow_rotating_primitives(fs, info);
      cfg.primitive_restart = info->primitive_restart;

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

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

      cfg.index_count = draw->count;
      cfg.index_type = panfrost_translate_index_size(info->index_size);

      if (PAN_ARCH >= 9) {
         /* Base vertex offset on Valhall is used for both
          * indexed and non-indexed draws, in a simple way for
          * either. Handle both cases.
          */
         if (cfg.index_type)
            cfg.base_vertex_offset = draw->index_bias;
         else
            cfg.base_vertex_offset = draw->start;

         /* Indices are moved outside the primitive descriptor
          * on Valhall, so we don't need to set that here
          */
      } else if (cfg.index_type) {
         cfg.base_vertex_offset = draw->index_bias - ctx->offset_start;

#if PAN_ARCH <= 7
         cfg.indices = batch->indices;
#endif
      }

#if PAN_ARCH >= 6
      cfg.secondary_shader = secondary_shader;
#endif
   }
}

#if PAN_ARCH == 9
static void
jm_emit_malloc_vertex_job(struct panfrost_batch *batch,
                          const struct pipe_draw_info *info,
                          const struct pipe_draw_start_count_bias *draw,
                          bool secondary_shader, void *job)
{
   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 fs_required = panfrost_fs_required(
      fs, ctx->blend, &ctx->pipe_framebuffer, ctx->depth_stencil);

   /* Varying shaders only feed data to the fragment shader, so if we omit
    * the fragment shader, we should omit the varying shader too.
    */
   secondary_shader &= fs_required;

   jm_emit_primitive(batch, info, draw, secondary_shader,
                     pan_section_ptr(job, MALLOC_VERTEX_JOB, PRIMITIVE));

   pan_section_pack(job, MALLOC_VERTEX_JOB, INSTANCE_COUNT, cfg) {
      cfg.count = info->instance_count;
   }

   pan_section_pack(job, MALLOC_VERTEX_JOB, ALLOCATION, cfg) {
      if (secondary_shader) {
         unsigned sz = panfrost_vertex_attribute_stride(vs, fs);
         cfg.vertex_packet_stride = sz + 16;
         cfg.vertex_attribute_stride = sz;
      } else {
         /* Hardware requirement for "no varyings" */
         cfg.vertex_packet_stride = 16;
         cfg.vertex_attribute_stride = 0;
      }
   }

   pan_section_pack(job, MALLOC_VERTEX_JOB, TILER, cfg) {
      cfg.address = jm_emit_tiler_desc(batch);
   }

   STATIC_ASSERT(sizeof(batch->scissor) == pan_size(SCISSOR));
   memcpy(pan_section_ptr(job, MALLOC_VERTEX_JOB, SCISSOR), &batch->scissor,
          pan_size(SCISSOR));

   panfrost_emit_primitive_size(
      ctx, info->mode == MESA_PRIM_POINTS, 0,
      pan_section_ptr(job, MALLOC_VERTEX_JOB, PRIMITIVE_SIZE));

   pan_section_pack(job, MALLOC_VERTEX_JOB, INDICES, cfg) {
      cfg.address = batch->indices;
   }

   jm_emit_tiler_draw(pan_section_ptr(job, MALLOC_VERTEX_JOB, DRAW), batch,
                      fs_required, u_reduced_prim(info->mode));

   pan_section_pack(job, MALLOC_VERTEX_JOB, POSITION, cfg) {
      jm_emit_shader_env(batch, &cfg, PIPE_SHADER_VERTEX,
                         panfrost_get_position_shader(batch, info));
   }

   pan_section_pack(job, MALLOC_VERTEX_JOB, VARYING, cfg) {
      /* If a varying shader is used, we configure it with the same
       * state as the position shader for backwards compatible
       * behaviour with Bifrost. This could be optimized.
       */
      if (!secondary_shader)
         continue;

      jm_emit_shader_env(batch, &cfg, PIPE_SHADER_VERTEX,
                         panfrost_get_varying_shader(batch));
   }
}
#endif

#if PAN_ARCH <= 7
static void
jm_emit_tiler_job(struct panfrost_batch *batch,
                  const struct pipe_draw_info *info,
                  const struct pipe_draw_start_count_bias *draw,
                  void *invocation_template, bool secondary_shader, void *job)
{
   struct panfrost_context *ctx = batch->ctx;

   void *section = pan_section_ptr(job, TILER_JOB, INVOCATION);
   memcpy(section, invocation_template, pan_size(INVOCATION));

   jm_emit_primitive(batch, info, draw, secondary_shader,
                     pan_section_ptr(job, TILER_JOB, PRIMITIVE));

   void *prim_size = pan_section_ptr(job, TILER_JOB, PRIMITIVE_SIZE);
   enum mesa_prim prim = u_reduced_prim(info->mode);

#if PAN_ARCH >= 6
   pan_section_pack(job, TILER_JOB, TILER, cfg) {
      cfg.address = jm_emit_tiler_desc(batch);
   }

   pan_section_pack(job, TILER_JOB, PADDING, cfg)
      ;
#endif

   jm_emit_tiler_draw(pan_section_ptr(job, TILER_JOB, DRAW), batch, true, prim);

   panfrost_emit_primitive_size(ctx, prim == MESA_PRIM_POINTS,
                                batch->varyings.psiz, prim_size);
}
#endif

void
GENX(jm_launch_xfb)(struct panfrost_batch *batch,
                    const struct pipe_draw_info *info, unsigned count)
{
   struct panfrost_ptr t = pan_pool_alloc_desc(&batch->pool.base, COMPUTE_JOB);

#if PAN_ARCH == 9
   pan_section_pack(t.cpu, COMPUTE_JOB, PAYLOAD, cfg) {
      cfg.workgroup_size_x = 1;
      cfg.workgroup_size_y = 1;
      cfg.workgroup_size_z = 1;

      cfg.workgroup_count_x = count;
      cfg.workgroup_count_y = info->instance_count;
      cfg.workgroup_count_z = 1;

      jm_emit_shader_env(batch, &cfg.compute, PIPE_SHADER_VERTEX,
                         batch->rsd[PIPE_SHADER_VERTEX]);

      /* TODO: Indexing. Also, this is a legacy feature... */
      cfg.compute.attribute_offset = batch->ctx->offset_start;

      /* Transform feedback shaders do not use barriers or shared
       * memory, so we may merge workgroups.
       */
      cfg.allow_merging_workgroups = true;
      cfg.task_increment = 1;
      cfg.task_axis = MALI_TASK_AXIS_Z;
   }
#else
   struct mali_invocation_packed invocation;

   panfrost_pack_work_groups_compute(&invocation, 1, count,
                                     info->instance_count, 1, 1, 1,
                                     PAN_ARCH <= 5, false);

   /* No varyings on XFB compute jobs. */
   mali_ptr saved_vs_varyings = batch->varyings.vs;

   batch->varyings.vs = 0;
   jm_emit_vertex_job(batch, info, &invocation, t.cpu);
   batch->varyings.vs = saved_vs_varyings;

#endif
   enum mali_job_type job_type = MALI_JOB_TYPE_COMPUTE;
#if PAN_ARCH <= 5
   job_type = MALI_JOB_TYPE_VERTEX;
#endif
   pan_jc_add_job(&batch->jm.jobs.vtc_jc, job_type, true, false, 0, 0, &t,
                  false);
}

#if PAN_ARCH < 9
/*
 * Push jobs required for the rasterization pipeline. If there are side effects
 * from the vertex shader, these are handled ahead-of-time with a compute
 * shader. This function should not be called if rasterization is skipped.
 */
static void
jm_push_vertex_tiler_jobs(struct panfrost_batch *batch,
                          const struct panfrost_ptr *vertex_job,
                          const struct panfrost_ptr *tiler_job)
{
   unsigned vertex =
      pan_jc_add_job(&batch->jm.jobs.vtc_jc, MALI_JOB_TYPE_VERTEX, false, false,
                     0, 0, vertex_job, false);

   pan_jc_add_job(&batch->jm.jobs.vtc_jc, MALI_JOB_TYPE_TILER, false, false,
                  vertex, 0, tiler_job, false);
}
#endif

void
GENX(jm_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 panfrost_context *ctx = batch->ctx;
   struct panfrost_compiled_shader *vs = ctx->prog[PIPE_SHADER_VERTEX];
   bool secondary_shader = vs->info.vs.secondary_enable;
   bool idvs = vs->info.vs.idvs;

#if PAN_ARCH <= 7
   struct mali_invocation_packed invocation;
   if (info->instance_count > 1) {
      panfrost_pack_work_groups_compute(&invocation, 1, vertex_count,
                                        info->instance_count, 1, 1, 1, true,
                                        false);
   } else {
      pan_pack(&invocation, INVOCATION, cfg) {
         cfg.invocations = vertex_count - 1;
         cfg.size_y_shift = 0;
         cfg.size_z_shift = 0;
         cfg.workgroups_x_shift = 0;
         cfg.workgroups_y_shift = 0;
         cfg.workgroups_z_shift = 32;
         cfg.thread_group_split = MALI_SPLIT_MIN_EFFICIENT;
      }
   }

   /* Emit all sort of descriptors. */
#endif

   UNUSED struct panfrost_ptr tiler, vertex;

   if (idvs) {
#if PAN_ARCH == 9
      tiler = pan_pool_alloc_desc(&batch->pool.base, MALLOC_VERTEX_JOB);
#elif PAN_ARCH >= 6
      tiler = pan_pool_alloc_desc(&batch->pool.base, INDEXED_VERTEX_JOB);
#else
      unreachable("IDVS is unsupported on Midgard");
#endif
   } else {
      vertex = pan_pool_alloc_desc(&batch->pool.base, COMPUTE_JOB);
      tiler = pan_pool_alloc_desc(&batch->pool.base, TILER_JOB);
   }

#if PAN_ARCH == 9
   assert(idvs && "Memory allocated IDVS required on Valhall");

   jm_emit_malloc_vertex_job(batch, info, draw, secondary_shader, tiler.cpu);

   pan_jc_add_job(&batch->jm.jobs.vtc_jc, MALI_JOB_TYPE_MALLOC_VERTEX, false,
                  false, 0, 0, &tiler, false);
#else
   /* Fire off the draw itself */
   jm_emit_tiler_job(batch, info, draw, &invocation, secondary_shader,
                     tiler.cpu);
   if (idvs) {
#if PAN_ARCH >= 6
      jm_emit_vertex_draw(
         batch, pan_section_ptr(tiler.cpu, INDEXED_VERTEX_JOB, VERTEX_DRAW));

      pan_jc_add_job(&batch->jm.jobs.vtc_jc, MALI_JOB_TYPE_INDEXED_VERTEX,
                     false, false, 0, 0, &tiler, false);
#endif
   } else {
      jm_emit_vertex_job(batch, info, &invocation, vertex.cpu);
      jm_push_vertex_tiler_jobs(batch, &vertex, &tiler);
   }
#endif
}

void
GENX(jm_launch_draw_indirect)(struct panfrost_batch *batch,
                              const struct pipe_draw_info *info,
                              unsigned drawid_offset,
                              const struct pipe_draw_indirect_info *indirect)
{
   unreachable("draw indirect not implemented for jm");
}

void
GENX(jm_emit_write_timestamp)(struct panfrost_batch *batch,
                              struct panfrost_resource *dst, unsigned offset)
{
   struct panfrost_ptr job =
      pan_pool_alloc_desc(&batch->pool.base, WRITE_VALUE_JOB);

   pan_section_pack(job.cpu, WRITE_VALUE_JOB, PAYLOAD, cfg) {
      cfg.address = dst->image.data.base + dst->image.data.offset + offset;
      cfg.type = MALI_WRITE_VALUE_TYPE_SYSTEM_TIMESTAMP;
   }

   pan_jc_add_job(&batch->jm.jobs.vtc_jc, MALI_JOB_TYPE_WRITE_VALUE, false,
                  false, 0, 0, &job, false);
   panfrost_batch_write_rsrc(batch, dst, PIPE_SHADER_VERTEX);
}