Indirect Draw Calls for deck.gl Instanced Layers

The exact sub-problem here is the per-frame attribute round-trip that deck.gl’s instanced layers perform when the visible set changes. A ScatterplotLayer or IconLayer rendering millions of instances re-derives its instance count and re-packs its per-instance attributes on the CPU whenever the viewport moves, then re-uploads the changed ranges before issuing a draw(vertexCount, instanceCount). During a continuous zoom or rotate that recomputation runs every frame: the CPU walks the feature array, tests each instance against the viewport, compacts the survivors into a fresh typed array, and pays a writeBuffer for the whole instance block — tens of megabytes for a dense point cloud — all before the GPU draws a single pixel. The instance count passed to draw() is a CPU-side number, so the GPU cannot begin until the CPU has finished culling. That serialization is the stall you see as a frame-time spike on the first frame of every interaction, and it scales with instance count, not with how many instances actually survive the cull.

An indirect draw breaks the round-trip by moving the count into GPU memory. Instead of draw(count) with a CPU-computed count, you call drawIndexedIndirect(buffer, offset), and the GPU reads the draw arguments — index count, instance count, and offsets — directly from a GPUBuffer. A compute pass populates that buffer: it culls the candidate instances, compacts the survivors’ ids with an atomicAdd cursor, and writes the surviving instance count into the args buffer’s second field. The CPU never learns the count, never re-packs attributes, and never uploads per-instance data after the initial load — it only updates a 16-byte viewport uniform. This page applies the compute-preprocessing pattern from the parent deck.gl layer integration with WebGPU reference specifically to the draw call itself, turning a CPU-gated instanced draw into a fully GPU-driven one. It assumes the compute versus render pipeline distinction and the persistent-buffer discipline that layer already established.

Runnable reference implementation

The indirect args buffer holds the five u32 fields drawIndexedIndirect expects — indexCount, instanceCount, firstIndex, baseVertex, firstInstance — in a 20-byte block. The four constant fields are seeded once at layer setup; only instanceCount changes per frame, and it is reset to zero on the CPU (a single 4-byte writeBuffer at offset 4) and then incremented on the GPU by the cull pass. The per-instance centers are uploaded exactly once. The buffer carries three usages at the same time: INDIRECT so the render pass can read draw args from it, STORAGE so the compute pass can write into it, and COPY_DST so the CPU can seed constants and reset the count.

// indirect-instances.ts — GPU-driven instanced draw for a deck.gl custom layer.
// drawIndexedIndirect args: [indexCount, instanceCount, firstIndex, baseVertex, firstInstance]
const INDIRECT_ARGS_SIZE = 5 * 4;        // 20 bytes, five u32
const INSTANCE_COUNT_OFFSET = 4;         // the second u32
const RESET_ZERO = new Uint32Array([0]); // reused reset payload, no per-frame alloc

interface IndirectResources {
  centerBuffer: GPUBuffer;    // per-instance vec2<f32> positions, uploaded ONCE
  visibleBuffer: GPUBuffer;   // compacted surviving instance ids, written by compute
  indirectBuffer: GPUBuffer;  // drawIndexedIndirect args, INDIRECT | STORAGE | COPY_DST
  viewportBuffer: GPUBuffer;  // uniform: the cull extent, the only per-frame CPU write
}

function createIndirectResources(
  device: GPUDevice,
  centers: Float32Array,      // packed vec2<f32>, one entry per candidate instance
  indicesPerInstance: number, // index count of the shared instance mesh
): IndirectResources {
  const instanceCount = centers.length / 2;

  const centerBuffer = device.createBuffer({
    label: "instance-centers",
    size: centers.byteLength,
    usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
  });
  device.queue.writeBuffer(centerBuffer, 0, centers);   // uploaded once, never per frame

  const visibleBuffer = device.createBuffer({
    label: "visible-instance-ids",
    size: instanceCount * 4,                             // worst case: all survive
    usage: GPUBufferUsage.STORAGE,
  });

  const indirectBuffer = device.createBuffer({
    label: "draw-args",
    size: INDIRECT_ARGS_SIZE,
    usage: GPUBufferUsage.INDIRECT | GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
  });
  // Seed constants once. instanceCount starts at 0; the cull pass fills it in.
  device.queue.writeBuffer(indirectBuffer, 0, new Uint32Array([
    indicesPerInstance, 0, 0, 0, 0,
  ]));

  const viewportBuffer = device.createBuffer({
    label: "cull-viewport",
    size: 16,                                            // vec2 min + vec2 max
    usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST,
  });

  return { centerBuffer, visibleBuffer, indirectBuffer, viewportBuffer };
}

The per-frame path updates only the viewport uniform and the instance-count reset, then encodes the cull compute pass and the indirect render pass into one command buffer. Queue operations are ordered, so the reset writeBuffer is guaranteed to land before the compute pass reads the args buffer, and the compute pass’s writes are visible to the render pass within the same submission.

function drawIndirectInstances(
  device: GPUDevice,
  res: IndirectResources,
  cull: { pipeline: GPUComputePipeline; bindGroup: GPUBindGroup },
  render: { pipeline: GPURenderPipeline; bindGroup: GPUBindGroup; indexBuffer: GPUBuffer },
  colorView: GPUTextureView,
  instanceCapacity: number,
  viewportExtent: Float32Array,   // [minX, minY, maxX, maxY] in projected space
): void {
  // The only per-frame CPU uploads: 16-byte extent + a 4-byte count reset.
  device.queue.writeBuffer(res.viewportBuffer, 0, viewportExtent);
  device.queue.writeBuffer(res.indirectBuffer, INSTANCE_COUNT_OFFSET, RESET_ZERO);

  const encoder = device.createCommandEncoder({ label: "indirect-frame" });

  // Compute: cull, compact ids into visibleBuffer, write instanceCount into the args.
  const pass = encoder.beginComputePass();
  pass.setPipeline(cull.pipeline);
  pass.setBindGroup(0, cull.bindGroup);
  pass.dispatchWorkgroups(Math.ceil(instanceCapacity / 64));
  pass.end();

  // Render: the GPU reads instanceCount straight from the buffer — no CPU readback.
  const rp = encoder.beginRenderPass({
    colorAttachments: [{
      view: colorView, loadOp: "clear", storeOp: "store",
      clearValue: { r: 0, g: 0, b: 0, a: 0 },
    }],
  });
  rp.setPipeline(render.pipeline);
  rp.setBindGroup(0, render.bindGroup);
  rp.setIndexBuffer(render.indexBuffer, "uint32");
  rp.drawIndexedIndirect(res.indirectBuffer, 0);   // draw args sourced from GPU memory
  rp.end();

  device.queue.submit([encoder.finish()]);
}

The compute kernel is what makes the count authoritative. It tests each candidate center against the viewport extent, and for every survivor claims a slot with atomicAdd on the args buffer’s instance_count field — the same atomic value doubles as the compaction cursor into visible. Declaring the args as a struct with an atomic<u32> in the instanceCount slot lets the render pass read the identical bytes as plain draw arguments; the atomic qualifier only governs how the compute pass writes them.

// cull.wgsl — populate the indirect args buffer and compact visible instances.
struct DrawArgs {
  index_count:    u32,           // constant, seeded by the CPU once
  instance_count: atomic<u32>,   // reset to 0 each frame, incremented here
  first_index:    u32,
  base_vertex:    i32,
  first_instance: u32,
};

struct Viewport { min: vec2<f32>, max: vec2<f32> };

@group(0) @binding(0) var<storage, read>       centers: array<vec2<f32>>; // candidates
@group(0) @binding(1) var<storage, read_write> visible: array<u32>;       // compacted ids
@group(0) @binding(2) var<storage, read_write> args:    DrawArgs;         // shared with the draw
@group(0) @binding(3) var<uniform>             view:    Viewport;

@compute @workgroup_size(64)
fn cull(@builtin(global_invocation_id) gid: vec3<u32>) {
  let i = gid.x;
  if (i >= arrayLength(&centers)) { return; }   // guard the ragged final workgroup

  let p = centers[i];
  // Viewport rejection in projected space; survivors contribute one instance.
  if (p.x < view.min.x || p.x > view.max.x ||
      p.y < view.min.y || p.y > view.max.y) {
    return;
  }
  // Claim a draw slot. instance_count is both the draw count and the write cursor.
  let slot = atomicAdd(&args.instance_count, 1u);
  visible[slot] = i;             // the vertex shader reads this to fetch attributes
}

The instance vertex shader indexes visible[instance_index] to recover the original instance id, then reads centerBuffer at that id — so the compacted survivor list drives attribute fetch entirely on the GPU. Because firstInstance stays zero, no optional feature is required; a non-zero firstInstance in an indirect draw needs the "indirect-first-instance" feature negotiated at device initialization for GIS workloads. The full bind-group wiring that joins the compute output to the render pass — matching @binding indices to deck.gl’s attribute slots so the framework does not force a recompile — is the subject of binding WebGPU render passes to deck.gl custom layers.

Parameter reference

Every structural value in the implementation above, with guidance for instanced spatial layers. These tables scroll horizontally on narrow viewports.

The reset of instanceCount before every compute pass is the constraint that bites first: omit it and the count accumulates across frames, so the second frame draws twice as many instances as survive, reading past the compacted region of visibleBuffer into stale ids. For the authoritative argument layout and alignment of drawIndexedIndirect, consult the W3C WebGPU specification and the MDN GPURenderPassEncoder.drawIndexedIndirect reference.

Failure modes

• Accumulating instance count. The instanceCount field is never reset, so it grows each frame and the draw reads past the compacted survivors. Detection: instance count climbs every frame in a captured trace; ghost or duplicated instances appear and the frame slows steadily. Fix: writeBuffer a zero to INSTANCE_COUNT_OFFSET every frame before the compute pass, as the reference does; queue ordering guarantees it lands first.
• Missing INDIRECT usage. The args buffer is created with STORAGE | COPY_DST only, and drawIndexedIndirect throws at draw time. Detection: GPUValidationError: indirect buffer usage ... does not contain INDIRECT. Fix: OR in GPUBufferUsage.INDIRECT; the buffer legitimately needs all three usages simultaneously.
• Under-sized visible buffer. visibleBuffer is sized to an average survivor count rather than the worst case, and a zoomed-out frame where every instance survives overflows the atomicAdd cursor. Detection: intermittent GPUValidationError on out-of-bounds storage writes, or garbage instances only at low zoom. Fix: size visibleBuffer to the full candidate count; the compaction can never exceed it.
• Race between reset and draw across submissions. Splitting the reset, compute, and draw across separate queue.submit() calls lets a second frame’s reset overtake the first frame’s draw, which reads a zeroed count. Detection: instances flicker to nothing on fast interaction. Fix: encode the reset dependency, compute pass, and indirect draw so the count is written and consumed within one ordered submission, and gate cross-frame reuse on queue.onSubmittedWorkDone(). To confirm the compute cull actually fits the frame budget, measure it with timestamp queries as covered in Performance Tuning & Profiling for WebGPU Spatial Pipelines and its frame profiling with timestamp queries guides.

Related

• deck.gl Layer Integration with WebGPU — the custom-layer lifecycle and compute-preprocessing pattern this draw path extends.
• Binding WebGPU Render Passes to deck.gl Custom Layers — wiring the compute output and indirect buffer into deck.gl’s render graph.
• WebGPU Compute vs Render Pipeline Fundamentals — why the cull belongs in a compute pass feeding the draw.
• Performance Tuning & Profiling for WebGPU Spatial Pipelines — measuring the frame-time win from removing the per-frame re-upload.

Up: deck.gl Layer Integration with WebGPU