// =====================================================================
// CC wedge size + prefix-sum compute shaders. Replaces the per-topology
// CPU loop in buildCcRenderState (cc_render.js:106) that walked
// runtimeIrregularInfo to fill wedgeNumDstPoints + wedgeStartPoints.
// The compute shaders here process EVERY loaded topology in a single
// dispatch each, regardless of how many meshes have been opened --
// not one shader run per mesh.
//
// Architecture:
//   - All wedges from all topologies are concatenated into one flat
//     input buffer (wedgeInBuf). meshTableBuf records each topology's
//     (wedgeBase, numWedges) slice.
//   - sizePipeline: one thread per wedge across the whole flat array,
//     computes numDstPoints = MAX_IRREG_LEVEL * (1 + valence*12 +
//     (border ? 3 : 0)). Writes wedgeOutBuf[i].x.
//   - prefixPipeline: one thread per MESH, serial-scans that mesh's
//     wedges (per-mesh exclusive prefix sum). Writes wedgeOutBuf[i].y.
//     Per-mesh scope keeps each topology's offsets starting at 0 so the
//     downstream wedge GPU dispatch + IrregularInfo blob layout don't
//     need to change.
//   - After the second pass, wedgeOutBuf is copied to a MAP_READ
//     buffer; refreshFromTopologies awaits the readback and writes the
//     GPU-computed sizes + offsets back into each ccState's
//     wedgeNumDstPoints + wedgeStartPoints (overwriting the prior CPU
//     values from buildCcRenderState). Caller verifies parity.
//
// Why not avoid the readback by having the wedge GPU shader / patch
// eval shader read offsets directly out of wedgeOutBuf at dispatch
// time? Future work. Today the CPU offsets are still read by
// appendCcStaticToTopoBlob's IrregularInfo packer + by CcGpuWedgeCtx
// .create's perWedgeCfg packer, so they have to live in ccState
// either way. The readback is a one-time cost per setCcScene call;
// the wedge offsets never change at runtime under MAX_IRREG_LEVEL
// pre-allocation.
// =====================================================================
'use strict';

const CC_GPU_WEDGE_SIZES_WGSL = /* wgsl */ `
struct WedgeIn { valence: u32, isBorder: u32, };
struct SizeParams { totalWedges: u32, maxLevel: u32, _pad0: u32, _pad1: u32, };

@group(0) @binding(0) var<storage, read>       wedgeIn  : array<WedgeIn>;
@group(0) @binding(1) var<storage, read_write> wedgeOut : array<vec2<u32>>;  // (numPts, start)
@group(0) @binding(2) var<uniform>             U        : SizeParams;

@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
  let i = gid.x;
  if (i >= U.totalWedges) { return; }
  let w = wedgeIn[i];
  var n = 1u + w.valence * 12u;
  if (w.isBorder != 0u) { n = n + 3u; }
  wedgeOut[i].x = U.maxLevel * n;
}
`;

// =====================================================================
// PER-FRAME three-pass variant (per-instance dynamic irregular level)
// =====================================================================
// Pass A_inst -- one thread per active instance. Each thread walks its
//   topology's wedges and accumulates the per-instance total size:
//     total[j] = sum_i (level[j] * (1 + valence_i*12 + (border_i?3:0)))
//
// Pass B_inst -- single-WG exclusive scan over per-instance totals to
//   produce the per-instance base offset into the suballocation:
//     base[j] = sum_{k<j} total[k]
//   Same role the existing per-mesh prefix played; one level up.
//
// Pass C_inst -- one thread per active instance. Each thread walks its
//   topology's wedges and writes the per-(instance, wedge) start offset
//   into the per-frame wedgeOffsetBuf:
//     offset[j,i] = base[j] + level[j] * exclusive_scan_within_instance(i)
//
// wedgeOffsetBuf is pre-allocated to maxInstances * maxWedgesPerTopology
// u32s and reused every frame. The per-(instance, wedge) addressing is
// flat-padded: index = j * maxWedgesPerTopology + i. Padded slots stay
// garbage but no consumer reads them; meshTable[topoIdx].numWedges
// bounds the legal range.
//
// Per-instance level lives in instInfoBuf as (topoIdx, level) pairs;
// numInstances is a uniform.
// =====================================================================

const CC_GPU_WEDGE_INST_TOTAL_WGSL = /* wgsl */ `
struct InstInfo  { topoIdx: u32, level: u32, };
struct MeshEntry { wedgeBase: u32, numWedges: u32, _pad0: u32, _pad1: u32, };
struct WedgeIn   { valence: u32, isBorder: u32, };
struct Params    { numInstances: u32, _pad0: u32, _pad1: u32, _pad2: u32, };

@group(0) @binding(0) var<storage, read>       instInfo  : array<InstInfo>;
@group(0) @binding(1) var<storage, read>       meshTable : array<MeshEntry>;
@group(0) @binding(2) var<storage, read>       wedgeIn   : array<WedgeIn>;
@group(0) @binding(3) var<storage, read_write> instTotal : array<u32>;
@group(0) @binding(4) var<uniform>             U         : Params;

@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
  let j = gid.x;
  if (j >= U.numInstances) { return; }
  let inst = instInfo[j];
  let m    = meshTable[inst.topoIdx];
  var total: u32 = 0u;
  for (var i: u32 = 0u; i < m.numWedges; i = i + 1u) {
    let w = wedgeIn[m.wedgeBase + i];
    var f = 1u + w.valence * 12u;
    if (w.isBorder != 0u) { f = f + 3u; }
    total = total + inst.level * f;
  }
  instTotal[j] = total;
}
`;

const CC_GPU_WEDGE_INST_PREFIX_WGSL = /* wgsl */ `
struct Params { numInstances: u32, _pad0: u32, _pad1: u32, _pad2: u32, };

@group(0) @binding(0) var<storage, read>       instTotal : array<u32>;
@group(0) @binding(1) var<storage, read_write> instBase  : array<u32>;
@group(0) @binding(2) var<uniform>             U         : Params;

// Serial single-WG exclusive scan, same shape as the per-mesh prefix
// shader above. With max_active_deformable_instances ~ 64, the cost is
// negligible; a multi-WG Blelloch scan is a follow-on if instance
// counts ever climb above a few hundred.
@compute @workgroup_size(1)
fn main() {
  var acc: u32 = 0u;
  for (var j: u32 = 0u; j < U.numInstances; j = j + 1u) {
    instBase[j] = acc;
    acc = acc + instTotal[j];
  }
}
`;

const CC_GPU_WEDGE_INST_LAYOUT_WGSL = /* wgsl */ `
struct InstInfo  { topoIdx: u32, level: u32, };
struct MeshEntry { wedgeBase: u32, numWedges: u32, _pad0: u32, _pad1: u32, };
struct WedgeIn   { valence: u32, isBorder: u32, };
struct Params    { numInstances: u32, maxWedgesPerTopology: u32, _pad0: u32, _pad1: u32, };

@group(0) @binding(0) var<storage, read>       instInfo    : array<InstInfo>;
@group(0) @binding(1) var<storage, read>       meshTable   : array<MeshEntry>;
@group(0) @binding(2) var<storage, read>       wedgeIn     : array<WedgeIn>;
@group(0) @binding(3) var<storage, read>       instBase    : array<u32>;
@group(0) @binding(4) var<storage, read_write> wedgeOffset : array<u32>;
@group(0) @binding(5) var<uniform>             U           : Params;

@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
  let j = gid.x;
  if (j >= U.numInstances) { return; }
  let inst = instInfo[j];
  let m    = meshTable[inst.topoIdx];
  let base = instBase[j];
  let row  = j * U.maxWedgesPerTopology;
  var acc: u32 = 0u;
  for (var i: u32 = 0u; i < m.numWedges; i = i + 1u) {
    let w = wedgeIn[m.wedgeBase + i];
    var f = 1u + w.valence * 12u;
    if (w.isBorder != 0u) { f = f + 3u; }
    wedgeOffset[row + i] = base + acc;
    acc = acc + inst.level * f;
  }
}
`;

const CC_GPU_WEDGE_PREFIX_WGSL = /* wgsl */ `
struct MeshEntry { wedgeBase: u32, numWedges: u32, _pad0: u32, _pad1: u32, };
struct PrefixParams { numMeshes: u32, _pad0: u32, _pad1: u32, _pad2: u32, };

@group(0) @binding(0) var<storage, read>       meshTable : array<MeshEntry>;
@group(0) @binding(1) var<storage, read_write> wedgeOut  : array<vec2<u32>>;
@group(0) @binding(2) var<uniform>             U         : PrefixParams;

// One thread per MESH; that thread serial-scans the mesh's slice of
// wedgeOut to write each wedge's exclusive prefix into .y. Serial
// scan is fine here because (a) this runs once per setCcScene call,
// not per frame, and (b) the largest fixture in the test set is imrod
// with ~7000 wedges, which is sub-millisecond at 1 GHz GPU clock.
// A multi-WG Blelloch scan would be the next step if per-frame
// runtime offset recompute becomes a feature (e.g. per-wedge LOD).
@compute @workgroup_size(1)
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
  let m = gid.x;
  if (m >= U.numMeshes) { return; }
  let e = meshTable[m];
  var acc : u32 = 0u;
  for (var i: u32 = 0u; i < e.numWedges; i = i + 1u) {
    let idx = e.wedgeBase + i;
    wedgeOut[idx].y = acc;
    acc = acc + wedgeOut[idx].x;
  }
}
`;

class CcGpuWedgeOffsetsCtx {
  static create(device, opts) {
    if (typeof device.createBuffer !== 'function') {
      throw new Error('CcGpuWedgeOffsetsCtx.create: device is not a GPUDevice');
    }
    const maxWedges            = (opts && opts.maxWedges) || 65536;
    const maxMeshes            = (opts && opts.maxMeshes) || 32;
    const MAX_LEVEL            = (opts && opts.maxLevel) || 8;
    // Per-frame variant caps. maxWedgesPerTopology is the per-instance
    // wedge-offset row width (flat-padded layout). Sized for the largest
    // shipping fixture (imrod ~7k wedges) with headroom.
    const maxInstances         = (opts && opts.maxInstances) || 64;
    const maxWedgesPerTopology = (opts && opts.maxWedgesPerTopology) || 8192;

    const wedgeInBuf = device.createBuffer({
      size:  Math.max(maxWedges * 8, 16),
      usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
      label: 'cc_gpu_offsets/wedgeIn',
    });
    const wedgeOutBuf = device.createBuffer({
      size:  Math.max(maxWedges * 8, 16),
      usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
      label: 'cc_gpu_offsets/wedgeOut',
    });
    const meshTableBuf = device.createBuffer({
      size:  Math.max(maxMeshes * 16, 16),
      usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
      label: 'cc_gpu_offsets/meshTable',
    });
    const sizeParamsBuf = device.createBuffer({
      size:  16,
      usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST,
      label: 'cc_gpu_offsets/sizeParams',
    });
    const prefixParamsBuf = device.createBuffer({
      size:  16,
      usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST,
      label: 'cc_gpu_offsets/prefixParams',
    });
    const readbackBuf = device.createBuffer({
      size:  Math.max(maxWedges * 8, 16),
      usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST,
      label: 'cc_gpu_offsets/readback',
    });

    // ----- Per-frame instance-level buffers ---------------------------
    const instInfoBuf = device.createBuffer({
      size:  Math.max(maxInstances * 8, 16),
      usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
      label: 'cc_gpu_offsets/instInfo',
    });
    const instTotalBuf = device.createBuffer({
      size:  Math.max(maxInstances * 4, 16),
      usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
      label: 'cc_gpu_offsets/instTotal',
    });
    const instBaseBuf = device.createBuffer({
      size:  Math.max(maxInstances * 4, 16),
      usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
      label: 'cc_gpu_offsets/instBase',
    });
    // Flat-padded per-(instance, wedge) start offsets into the wedge
    // eval suballocation. Bound by downstream wedge eval + patch eval.
    const wedgeOffsetBuf = device.createBuffer({
      size:  Math.max(maxInstances * maxWedgesPerTopology * 4, 16),
      usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
      label: 'cc_gpu_offsets/wedgeOffset',
    });
    const instTotalParamsBuf = device.createBuffer({
      size:  16,
      usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST,
      label: 'cc_gpu_offsets/instTotalParams',
    });
    const instPrefixParamsBuf = device.createBuffer({
      size:  16,
      usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST,
      label: 'cc_gpu_offsets/instPrefixParams',
    });
    const instLayoutParamsBuf = device.createBuffer({
      size:  16,
      usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST,
      label: 'cc_gpu_offsets/instLayoutParams',
    });

    const sizeLayout = device.createBindGroupLayout({
      entries: [
        { binding: 0, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'read-only-storage' } },
        { binding: 1, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'storage' } },
        { binding: 2, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'uniform' } },
      ],
    });
    const sizePipeline = device.createComputePipeline({
      layout:  device.createPipelineLayout({ bindGroupLayouts: [sizeLayout] }),
      compute: { module: device.createShaderModule({ code: CC_GPU_WEDGE_SIZES_WGSL }), entryPoint: 'main' },
    });
    const sizeBindGroup = device.createBindGroup({
      layout: sizeLayout,
      entries: [
        { binding: 0, resource: { buffer: wedgeInBuf       } },
        { binding: 1, resource: { buffer: wedgeOutBuf      } },
        { binding: 2, resource: { buffer: sizeParamsBuf    } },
      ],
    });

    const prefixLayout = device.createBindGroupLayout({
      entries: [
        { binding: 0, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'read-only-storage' } },
        { binding: 1, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'storage' } },
        { binding: 2, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'uniform' } },
      ],
    });
    const prefixPipeline = device.createComputePipeline({
      layout:  device.createPipelineLayout({ bindGroupLayouts: [prefixLayout] }),
      compute: { module: device.createShaderModule({ code: CC_GPU_WEDGE_PREFIX_WGSL }), entryPoint: 'main' },
    });
    const prefixBindGroup = device.createBindGroup({
      layout: prefixLayout,
      entries: [
        { binding: 0, resource: { buffer: meshTableBuf     } },
        { binding: 1, resource: { buffer: wedgeOutBuf      } },
        { binding: 2, resource: { buffer: prefixParamsBuf  } },
      ],
    });

    // ----- Per-frame Pass A_inst (per-instance total) -----------------
    const instTotalLayout = device.createBindGroupLayout({
      entries: [
        { binding: 0, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'read-only-storage' } },
        { binding: 1, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'read-only-storage' } },
        { binding: 2, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'read-only-storage' } },
        { binding: 3, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'storage' } },
        { binding: 4, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'uniform' } },
      ],
    });
    const instTotalPipeline = device.createComputePipeline({
      layout:  device.createPipelineLayout({ bindGroupLayouts: [instTotalLayout] }),
      compute: { module: device.createShaderModule({ code: CC_GPU_WEDGE_INST_TOTAL_WGSL }), entryPoint: 'main' },
    });
    const instTotalBindGroup = device.createBindGroup({
      layout: instTotalLayout,
      entries: [
        { binding: 0, resource: { buffer: instInfoBuf        } },
        { binding: 1, resource: { buffer: meshTableBuf       } },
        { binding: 2, resource: { buffer: wedgeInBuf         } },
        { binding: 3, resource: { buffer: instTotalBuf       } },
        { binding: 4, resource: { buffer: instTotalParamsBuf } },
      ],
    });

    // ----- Per-frame Pass B_inst (exclusive scan over totals) ---------
    const instPrefixLayout = device.createBindGroupLayout({
      entries: [
        { binding: 0, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'read-only-storage' } },
        { binding: 1, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'storage' } },
        { binding: 2, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'uniform' } },
      ],
    });
    const instPrefixPipeline = device.createComputePipeline({
      layout:  device.createPipelineLayout({ bindGroupLayouts: [instPrefixLayout] }),
      compute: { module: device.createShaderModule({ code: CC_GPU_WEDGE_INST_PREFIX_WGSL }), entryPoint: 'main' },
    });
    const instPrefixBindGroup = device.createBindGroup({
      layout: instPrefixLayout,
      entries: [
        { binding: 0, resource: { buffer: instTotalBuf        } },
        { binding: 1, resource: { buffer: instBaseBuf         } },
        { binding: 2, resource: { buffer: instPrefixParamsBuf } },
      ],
    });

    // ----- Per-frame Pass C_inst (per-(inst, wedge) layout) -----------
    const instLayoutLayout = device.createBindGroupLayout({
      entries: [
        { binding: 0, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'read-only-storage' } },
        { binding: 1, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'read-only-storage' } },
        { binding: 2, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'read-only-storage' } },
        { binding: 3, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'read-only-storage' } },
        { binding: 4, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'storage' } },
        { binding: 5, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'uniform' } },
      ],
    });
    const instLayoutPipeline = device.createComputePipeline({
      layout:  device.createPipelineLayout({ bindGroupLayouts: [instLayoutLayout] }),
      compute: { module: device.createShaderModule({ code: CC_GPU_WEDGE_INST_LAYOUT_WGSL }), entryPoint: 'main' },
    });
    const instLayoutBindGroup = device.createBindGroup({
      layout: instLayoutLayout,
      entries: [
        { binding: 0, resource: { buffer: instInfoBuf         } },
        { binding: 1, resource: { buffer: meshTableBuf        } },
        { binding: 2, resource: { buffer: wedgeInBuf          } },
        { binding: 3, resource: { buffer: instBaseBuf         } },
        { binding: 4, resource: { buffer: wedgeOffsetBuf      } },
        { binding: 5, resource: { buffer: instLayoutParamsBuf } },
      ],
    });

    return new CcGpuWedgeOffsetsCtx({
      device, maxWedges, maxMeshes, MAX_LEVEL,
      maxInstances, maxWedgesPerTopology,
      wedgeInBuf, wedgeOutBuf, meshTableBuf,
      sizeParamsBuf, prefixParamsBuf, readbackBuf,
      sizePipeline, prefixPipeline,
      sizeBindGroup, prefixBindGroup,
      // Per-frame fields
      instInfoBuf, instTotalBuf, instBaseBuf, wedgeOffsetBuf,
      instTotalParamsBuf, instPrefixParamsBuf, instLayoutParamsBuf,
      instTotalPipeline, instPrefixPipeline, instLayoutPipeline,
      instTotalBindGroup, instPrefixBindGroup, instLayoutBindGroup,
      _busy: false,
    });
  }

  constructor(fields) { Object.assign(this, fields); }

  // ONE size dispatch + ONE prefix-sum dispatch covering every loaded
  // topology. Returns { meshes, totalWedges, maxDeltaVsCpu } so callers
  // can gate on bit-exact parity with the CPU values that were already
  // in ccState.wedgeStartPoints / wedgeNumDstPoints.
  //
  // After this resolves, every loaded topology's ccState carries the
  // GPU-computed sizes + offsets (which match the CPU values bit-for-
  // bit by construction -- same formula, same MAX_LEVEL).
  async refreshFromTopologies(ccTopologies, pendingName, pendingCcState) {
    if (this._busy) {
      throw new Error('CcGpuWedgeOffsetsCtx.refreshFromTopologies: already in progress (concurrent calls not supported)');
    }
    this._busy = true;
    try {
      const entries = [];
      const seenNames = new Set();
      function push(name, ccs) {
        if (!ccs || seenNames.has(name)) return;
        if (!ccs.patches || !ccs.patches.m_weightInflData ||
            !ccs.patches.m_weightInflData.m_runtimeIrregularInfo) return;
        seenNames.add(name);
        entries.push({ name, ccs });
      }
      // Walk ccTopologies, but substitute pendingCcState whenever its
      // name matches an existing entry. This handles the common case
      // where setCcScene is re-loading a fixture that already lives
      // in ccTopologies -- the cached topology still carries the OLD
      // ccState while the in-flight ccState (just produced by
      // buildCcRenderState) is what the rest of setCcScene is about
      // to install. Without the substitution the refresh would emit
      // BOTH ccStates as separate "meshes" and double-count.
      const cacheName = pendingName || '<pending>';
      if (pendingCcState) push(cacheName, pendingCcState);
      for (const k of Object.keys(ccTopologies)) {
        const t = ccTopologies[k];
        if (t) push(k, t.ccState);
      }
      if (entries.length === 0) return { meshes: 0, totalWedges: 0, maxDeltaVsCpu: 0 };
      if (entries.length > this.maxMeshes) {
        throw new Error('CcGpuWedgeOffsetsCtx: ' + entries.length + ' meshes exceeds capacity ' + this.maxMeshes);
      }

      // Pack one flat wedge-input array + one mesh table covering ALL
      // entries.
      let totalWedges = 0;
      const ranges = [];
      for (let m = 0; m < entries.length; m++) {
        const wi = entries[m].ccs.patches.m_weightInflData.m_runtimeIrregularInfo;
        ranges.push({ base: totalWedges, count: wi.length });
        totalWedges += wi.length;
      }
      if (totalWedges === 0) return { meshes: entries.length, totalWedges: 0, maxDeltaVsCpu: 0 };
      if (totalWedges > this.maxWedges) {
        throw new Error('CcGpuWedgeOffsetsCtx: ' + totalWedges + ' wedges exceeds capacity ' + this.maxWedges);
      }
      const wedgeInArr   = new Uint32Array(Math.max(totalWedges * 2, 4));
      const meshTableArr = new Uint32Array(Math.max(entries.length * 4, 4));
      for (let m = 0; m < entries.length; m++) {
        const r = ranges[m];
        meshTableArr[m * 4 + 0] = r.base;
        meshTableArr[m * 4 + 1] = r.count;
        const wi = entries[m].ccs.patches.m_weightInflData.m_runtimeIrregularInfo;
        for (let i = 0; i < r.count; i++) {
          const w = wi[i];
          wedgeInArr[(r.base + i) * 2 + 0] = w.m_valence >>> 0;
          wedgeInArr[(r.base + i) * 2 + 1] = w.m_isBorder ? 1 : 0;
        }
      }

      this.device.queue.writeBuffer(this.wedgeInBuf, 0,
        wedgeInArr.buffer, wedgeInArr.byteOffset, wedgeInArr.byteLength);
      this.device.queue.writeBuffer(this.meshTableBuf, 0,
        meshTableArr.buffer, meshTableArr.byteOffset, meshTableArr.byteLength);
      this.device.queue.writeBuffer(this.sizeParamsBuf, 0,
        new Uint32Array([totalWedges, this.MAX_LEVEL, 0, 0]).buffer);
      this.device.queue.writeBuffer(this.prefixParamsBuf, 0,
        new Uint32Array([entries.length, 0, 0, 0]).buffer);

      const enc = this.device.createCommandEncoder({ label: 'cc_gpu_offsets/refresh' });
      const sizePass = enc.beginComputePass({ label: 'cc_gpu_offsets/size' });
      sizePass.setPipeline(this.sizePipeline);
      sizePass.setBindGroup(0, this.sizeBindGroup);
      sizePass.dispatchWorkgroups(Math.ceil(totalWedges / 64));
      sizePass.end();
      const prefixPass = enc.beginComputePass({ label: 'cc_gpu_offsets/prefix' });
      prefixPass.setPipeline(this.prefixPipeline);
      prefixPass.setBindGroup(0, this.prefixBindGroup);
      prefixPass.dispatchWorkgroups(entries.length);
      prefixPass.end();
      const readbackBytes = totalWedges * 8;
      enc.copyBufferToBuffer(this.wedgeOutBuf, 0, this.readbackBuf, 0, readbackBytes);
      this.device.queue.submit([enc.finish()]);

      await this.readbackBuf.mapAsync(GPUMapMode.READ, 0, readbackBytes);
      const result = new Uint32Array(this.readbackBuf.getMappedRange(0, readbackBytes).slice(0));
      this.readbackBuf.unmap();

      // Write GPU offsets back into each topology's ccState. Compare
      // to whatever CPU values were sitting there.
      let maxDelta = 0;
      for (let m = 0; m < entries.length; m++) {
        const r = ranges[m];
        const ccs = entries[m].ccs;
        for (let i = 0; i < r.count; i++) {
          const gpuNum   = result[(r.base + i) * 2 + 0];
          const gpuStart = result[(r.base + i) * 2 + 1];
          const cpuNum   = (ccs.wedgeNumDstPoints && ccs.wedgeNumDstPoints[i]) | 0;
          const cpuStart = (ccs.wedgeStartPoints  && ccs.wedgeStartPoints [i]) | 0;
          if (Math.abs(gpuNum   - cpuNum)   > maxDelta) maxDelta = Math.abs(gpuNum   - cpuNum);
          if (Math.abs(gpuStart - cpuStart) > maxDelta) maxDelta = Math.abs(gpuStart - cpuStart);
          if (ccs.wedgeNumDstPoints) ccs.wedgeNumDstPoints[i] = gpuNum;
          if (ccs.wedgeStartPoints)  ccs.wedgeStartPoints [i] = gpuStart;
        }
      }
      return { meshes: entries.length, totalWedges, maxDeltaVsCpu: maxDelta };
    } finally {
      this._busy = false;
    }
  }

  // -------------------------------------------------------------------
  // Per-frame three-pass dispatch.
  //
  // Encodes Pass A_inst, Pass B_inst, Pass C_inst into the caller's
  // command encoder and returns. The caller is responsible for
  // submitting the encoder; the per-frame buffers (wedgeOffsetBuf,
  // instBaseBuf, instTotalBuf) are populated by the time the encoder
  // is submitted and any downstream consumer reads them.
  //
  // Preconditions:
  //   - refreshFromTopologies (or its equivalent) has already written
  //     the static wedgeIn + meshTable for all topologies that the
  //     instances reference.
  //   - instInfoArr is an array of { topoIdx, level } entries with
  //     length <= maxInstances.
  //
  // Buffer layouts after dispatch:
  //   instTotalBuf[j]   = per-instance total size (u32)
  //   instBaseBuf[j]    = exclusive prefix-sum of instTotal (u32)
  //   wedgeOffsetBuf[j * maxWedgesPerTopology + i]
  //                     = base offset for instance j, wedge i into the
  //                       wedge eval suballocation (u32)
  //
  // No readback, no CPU sync. Designed to run every frame from the
  // engine's render loop, before the irregular wedge eval dispatch.
  // -------------------------------------------------------------------
  dispatchPerFrame(encoder, instInfoArr) {
    if (!Array.isArray(instInfoArr)) {
      throw new Error('CcGpuWedgeOffsetsCtx.dispatchPerFrame: instInfoArr must be an array');
    }
    const numInstances = instInfoArr.length;
    if (numInstances === 0) return { numInstances: 0 };
    if (numInstances > this.maxInstances) {
      throw new Error('CcGpuWedgeOffsetsCtx.dispatchPerFrame: ' + numInstances +
                      ' instances exceeds capacity ' + this.maxInstances);
    }

    // Pack instInfo: u32 topoIdx + u32 level per instance.
    const instArr = new Uint32Array(Math.max(numInstances * 2, 4));
    for (let j = 0; j < numInstances; j++) {
      const e = instInfoArr[j];
      const topoIdx = (e.topoIdx | 0) >>> 0;
      const level   = (e.level   | 0) >>> 0;
      if (topoIdx >= this.maxMeshes) {
        throw new Error('CcGpuWedgeOffsetsCtx.dispatchPerFrame: instance ' + j +
                        ' has topoIdx=' + topoIdx + ' >= maxMeshes=' + this.maxMeshes);
      }
      if (level === 0 || level > this.MAX_LEVEL) {
        throw new Error('CcGpuWedgeOffsetsCtx.dispatchPerFrame: instance ' + j +
                        ' level=' + level + ' must be in [1, ' + this.MAX_LEVEL + ']');
      }
      instArr[j * 2 + 0] = topoIdx;
      instArr[j * 2 + 1] = level;
    }
    this.device.queue.writeBuffer(this.instInfoBuf, 0,
      instArr.buffer, instArr.byteOffset, instArr.byteLength);
    this.device.queue.writeBuffer(this.instTotalParamsBuf, 0,
      new Uint32Array([numInstances, 0, 0, 0]).buffer);
    this.device.queue.writeBuffer(this.instPrefixParamsBuf, 0,
      new Uint32Array([numInstances, 0, 0, 0]).buffer);
    this.device.queue.writeBuffer(this.instLayoutParamsBuf, 0,
      new Uint32Array([numInstances, this.maxWedgesPerTopology, 0, 0]).buffer);

    const totalPass = encoder.beginComputePass({ label: 'cc_gpu_offsets/inst_total' });
    totalPass.setPipeline(this.instTotalPipeline);
    totalPass.setBindGroup(0, this.instTotalBindGroup);
    totalPass.dispatchWorkgroups(Math.ceil(numInstances / 64));
    totalPass.end();

    const prefixPass = encoder.beginComputePass({ label: 'cc_gpu_offsets/inst_prefix' });
    prefixPass.setPipeline(this.instPrefixPipeline);
    prefixPass.setBindGroup(0, this.instPrefixBindGroup);
    prefixPass.dispatchWorkgroups(1);
    prefixPass.end();

    const layoutPass = encoder.beginComputePass({ label: 'cc_gpu_offsets/inst_layout' });
    layoutPass.setPipeline(this.instLayoutPipeline);
    layoutPass.setBindGroup(0, this.instLayoutBindGroup);
    layoutPass.dispatchWorkgroups(Math.ceil(numInstances / 64));
    layoutPass.end();

    return { numInstances };
  }

  destroy() {
    this.wedgeInBuf.destroy();
    this.wedgeOutBuf.destroy();
    this.meshTableBuf.destroy();
    this.sizeParamsBuf.destroy();
    this.prefixParamsBuf.destroy();
    this.readbackBuf.destroy();
    this.instInfoBuf.destroy();
    this.instTotalBuf.destroy();
    this.instBaseBuf.destroy();
    this.wedgeOffsetBuf.destroy();
    this.instTotalParamsBuf.destroy();
    this.instPrefixParamsBuf.destroy();
    this.instLayoutParamsBuf.destroy();
  }
}

if (typeof module !== 'undefined' && module.exports) {
  module.exports = { CcGpuWedgeOffsetsCtx,
                     CC_GPU_WEDGE_SIZES_WGSL,
                     CC_GPU_WEDGE_PREFIX_WGSL,
                     CC_GPU_WEDGE_INST_TOTAL_WGSL,
                     CC_GPU_WEDGE_INST_PREFIX_WGSL,
                     CC_GPU_WEDGE_INST_LAYOUT_WGSL };
}
