// =====================================================================
// CC compute-shader port (Phase 6a+). WGSL ports of the per-frame
// stencil application. Standalone of the rest of the viewer's
// tessellation pipeline -- each function builds its own pipelines
// + bind groups, uploads the relevant buffers, dispatches, reads
// back, and returns a typed-array result.
//
// Validation pattern: the in-browser harness calls
// `gpuComputeCcLimitPositions(device, ccState, srcPositions)` and
// diffs the returned positions against the CPU-side `state.limitPosData`
// already produced by `ccBeginFrame`. CPU is the gold reference
// (already byte-equal vs amolite per Phase 2's CC_LimitPositions gate),
// so a small tolerance (~1e-4 vs f32 GPU rounding) is fine.
'use strict';

// =====================================================================
// CC_GPU_LIMIT_PN_WGSL
//
// One workgroup thread per output limit vertex. Reads VertexLimitInfo's
// (dataStart, dataNum) + (influenceData, weightDataP, weightDataT0,
// weightDataT1), accumulates three weighted stencils over source
// positions to get (P, T0, T1), then N = normalize(cross(T0, T1)).
// Writes P + N to two output buffers.
//
// Mirrors Phase 2's CalcLimitPointsAndNormals (CPU) byte-equal modulo
// f32 reassociation. The single-pass design folds limit-positions
// (Phase 6a) and limit-normals (Phase 6b) into one workgroup dispatch
// to minimize WebGPU command-encoder cost per frame.
//
// Layout:
//   limitDataStart[v]  : i32 -- offset into influenceData/weightData*
//   limitDataNum[v]    : i32 -- number of influences for vertex v
//   limitInfluenceData[]: i32 (one influence per slot)
//   limitWeightDataP[] : f32 (one position weight per slot)
//   limitWeightDataT0[]: f32 (one tangent-0 weight per slot)
//   limitWeightDataT1[]: f32 (one tangent-1 weight per slot)
//   srcPositions[]     : vec3<f32> packed flat
//   outLimitPos[]      : vec3<f32> packed flat
//   outLimitNrm[]      : vec3<f32> packed flat
//
// (Future Phase 6c: irregular wedge subdiv compute. 6d: patch eval
// inside emit_verts.)
// =====================================================================
const CC_GPU_LIMIT_PN_WGSL = /* wgsl */ `
struct Params {
  numLimitPoints : u32,
  lpOffWords     : u32,  // f32-word offset where limit positions land in outBuf
  lnOffWords     : u32,  // f32-word offset where limit normals land in outBuf
  _pad0          : u32,
};
@group(0) @binding(0) var<uniform>                P             : Params;
@group(0) @binding(1) var<storage, read>          dataStart     : array<i32>;
@group(0) @binding(2) var<storage, read>          dataNum       : array<i32>;
@group(0) @binding(3) var<storage, read>          influenceData : array<i32>;
@group(0) @binding(4) var<storage, read>          weightDataP   : array<f32>;
@group(0) @binding(5) var<storage, read>          weightDataT0  : array<f32>;
@group(0) @binding(6) var<storage, read>          weightDataT1  : array<f32>;
@group(0) @binding(7) var<storage, read>          srcPositions  : array<f32>;  // xyz flat
// Single output buffer; the runtime path binds topo.gpuBlob directly here
// and the shader writes pos at P.lpOffWords and nrm at P.lnOffWords (both
// in f32-word units). Eliminates the intermediate outPosBuf/outNrmBuf and
// the two copyBufferToBuffer calls that used to land them into gpuBlob --
// those copies' UAV->COPY_SRC transition barriers were being attributed to
// limit_pn's end-of-pass timestamp (see investigation in chat).
@group(0) @binding(8) var<storage, read_write>    outBuf        : array<f32>;

@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid : vec3<u32>) {
  let v = gid.x;
  if (v >= P.numLimitPoints) { return; }
  let start = u32(dataStart[v]);
  let num   = u32(dataNum[v]);

  var pos  = vec3<f32>(0.0);
  var tan0 = vec3<f32>(0.0);
  var tan1 = vec3<f32>(0.0);

  for (var k : u32 = 0u; k < num; k = k + 1u) {
    let infl = u32(influenceData[start + k]);
    let wp   = weightDataP [start + k];
    let wt0  = weightDataT0[start + k];
    let wt1  = weightDataT1[start + k];
    let s = vec3<f32>(
      srcPositions[infl * 3u + 0u],
      srcPositions[infl * 3u + 1u],
      srcPositions[infl * 3u + 2u]);
    pos  = pos  + s * wp;
    tan0 = tan0 + s * wt0;
    tan1 = tan1 + s * wt1;
  }

  let nRaw = cross(tan0, tan1);
  let nLen = length(nRaw);
  let nrm  = select(vec3<f32>(0.0, 0.0, 0.0), nRaw / nLen, nLen > 0.0);

  outBuf[P.lpOffWords + v * 3u + 0u] = pos.x;
  outBuf[P.lpOffWords + v * 3u + 1u] = pos.y;
  outBuf[P.lpOffWords + v * 3u + 2u] = pos.z;
  outBuf[P.lnOffWords + v * 3u + 0u] = nrm.x;
  outBuf[P.lnOffWords + v * 3u + 1u] = nrm.y;
  outBuf[P.lnOffWords + v * 3u + 2u] = nrm.z;
}
`;

// Legacy alias for callers that still reference the Phase 6a-only
// shader name. Internally we always use the PN combined pass.
const CC_GPU_LIMIT_POS_WGSL = CC_GPU_LIMIT_PN_WGSL;

// ---- Variant A: workgroup_size(32) ----
// Same per-thread work as the default, just a smaller workgroup. The
// SOLE point is dispatch breadth: the default's workgroup_size(64) for
// imrod's 8179 limit points = 128 WGs / 46 SMs ~= 2.8 WGs/SM, far
// below the SM's 16-WG resident cap, so a stall on global-mem latency
// has no other warp to swap in. workgroup_size(32) doubles the WG
// count to 256 (= 5.5 WGs/SM) without changing per-thread work or
// register pressure -- cheap latency-hiding test.
const CC_GPU_LIMIT_PN_WS32_WGSL = /* wgsl */ `
struct Params {
  numLimitPoints : u32,
  lpOffWords     : u32,
  lnOffWords     : u32,
  _pad0          : u32,
};
@group(0) @binding(0) var<uniform>             P             : Params;
@group(0) @binding(1) var<storage, read>       dataStart     : array<i32>;
@group(0) @binding(2) var<storage, read>       dataNum       : array<i32>;
@group(0) @binding(3) var<storage, read>       influenceData : array<i32>;
@group(0) @binding(4) var<storage, read>       weightDataP   : array<f32>;
@group(0) @binding(5) var<storage, read>       weightDataT0  : array<f32>;
@group(0) @binding(6) var<storage, read>       weightDataT1  : array<f32>;
@group(0) @binding(7) var<storage, read>       srcPositions  : array<f32>;
@group(0) @binding(8) var<storage, read_write> outBuf        : array<f32>;

@compute @workgroup_size(32)
fn main(@builtin(global_invocation_id) gid : vec3<u32>) {
  let v = gid.x;
  if (v >= P.numLimitPoints) { return; }
  let start = u32(dataStart[v]);
  let num   = u32(dataNum[v]);
  var pos  = vec3<f32>(0.0);
  var tan0 = vec3<f32>(0.0);
  var tan1 = vec3<f32>(0.0);
  for (var k : u32 = 0u; k < num; k = k + 1u) {
    let infl = u32(influenceData[start + k]);
    let s = vec3<f32>(
      srcPositions[infl * 3u + 0u],
      srcPositions[infl * 3u + 1u],
      srcPositions[infl * 3u + 2u]);
    pos  = pos  + s * weightDataP [start + k];
    tan0 = tan0 + s * weightDataT0[start + k];
    tan1 = tan1 + s * weightDataT1[start + k];
  }
  let nRaw = cross(tan0, tan1);
  let nLen = length(nRaw);
  let nrm  = select(vec3<f32>(0.0, 0.0, 0.0), nRaw / nLen, nLen > 0.0);
  outBuf[P.lpOffWords + v * 3u + 0u] = pos.x;
  outBuf[P.lpOffWords + v * 3u + 1u] = pos.y;
  outBuf[P.lpOffWords + v * 3u + 2u] = pos.z;
  outBuf[P.lnOffWords + v * 3u + 0u] = nrm.x;
  outBuf[P.lnOffWords + v * 3u + 1u] = nrm.y;
  outBuf[P.lnOffWords + v * 3u + 2u] = nrm.z;
}
`;

// ---- Variant B: cooperative-per-vertex (8 lanes per limit point) ----
// workgroup_size stays at 64 but now 8 threads cooperate on one limit
// point: 8 verts per WG, each vert's 8 lanes split the influence loop
// strided by 8. Partial pos/tan0/tan1 land in workgroup shared mem,
// reduced by lane 0 of each vert group with a single workgroupBarrier.
//
// Wins on two axes:
//   (1) dispatch breadth -- imrod's 128 WGs become 8179/8 ~= 1023 WGs,
//       fully saturating the SM cap (16 WGs/SM x 46 SMs = 736 slots),
//       so the SM scheduler always has work to swap in on a global-
//       mem stall.
//   (2) inner-loop length -- avg per-thread loop count drops from 10.5
//       to 1.3 (max from 21 to 3). Warp-divergence tail loss largely
//       goes away.
//
// Cost: 1 workgroupBarrier per dispatch, ~3 KB workgroup mem per WG.
// Total threads dispatched grows by 8x but the SM is now actually
// using them; the alternative was idle warps.
const CC_GPU_LIMIT_PN_COOP_WGSL = /* wgsl */ `
struct Params {
  numLimitPoints : u32,
  lpOffWords     : u32,
  lnOffWords     : u32,
  _pad0          : u32,
};
@group(0) @binding(0) var<uniform>             P             : Params;
@group(0) @binding(1) var<storage, read>       dataStart     : array<i32>;
@group(0) @binding(2) var<storage, read>       dataNum       : array<i32>;
@group(0) @binding(3) var<storage, read>       influenceData : array<i32>;
@group(0) @binding(4) var<storage, read>       weightDataP   : array<f32>;
@group(0) @binding(5) var<storage, read>       weightDataT0  : array<f32>;
@group(0) @binding(6) var<storage, read>       weightDataT1  : array<f32>;
@group(0) @binding(7) var<storage, read>       srcPositions  : array<f32>;
@group(0) @binding(8) var<storage, read_write> outBuf        : array<f32>;

const LANES_PER_VERT : u32 = 8u;
const VERTS_PER_WG   : u32 = 8u;   // 64 / 8

var<workgroup> shPos  : array<vec3<f32>, 64>;
var<workgroup> shTan0 : array<vec3<f32>, 64>;
var<workgroup> shTan1 : array<vec3<f32>, 64>;

@compute @workgroup_size(64)
fn main(@builtin(workgroup_id) wid : vec3<u32>,
        @builtin(local_invocation_id) lid : vec3<u32>) {
  let vertInWg   = lid.x / LANES_PER_VERT;
  let laneInVert = lid.x % LANES_PER_VERT;
  let v          = wid.x * VERTS_PER_WG + vertInWg;

  var pos  = vec3<f32>(0.0);
  var tan0 = vec3<f32>(0.0);
  var tan1 = vec3<f32>(0.0);
  if (v < P.numLimitPoints) {
    let start = u32(dataStart[v]);
    let num   = u32(dataNum[v]);
    for (var k : u32 = laneInVert; k < num; k = k + LANES_PER_VERT) {
      let infl = u32(influenceData[start + k]);
      let s = vec3<f32>(
        srcPositions[infl * 3u + 0u],
        srcPositions[infl * 3u + 1u],
        srcPositions[infl * 3u + 2u]);
      pos  = pos  + s * weightDataP [start + k];
      tan0 = tan0 + s * weightDataT0[start + k];
      tan1 = tan1 + s * weightDataT1[start + k];
    }
  }
  shPos [lid.x] = pos;
  shTan0[lid.x] = tan0;
  shTan1[lid.x] = tan1;
  workgroupBarrier();

  if (laneInVert == 0u && v < P.numLimitPoints) {
    var p  = vec3<f32>(0.0);
    var t0 = vec3<f32>(0.0);
    var t1 = vec3<f32>(0.0);
    let base = vertInWg * LANES_PER_VERT;
    for (var i : u32 = 0u; i < LANES_PER_VERT; i = i + 1u) {
      p  = p  + shPos [base + i];
      t0 = t0 + shTan0[base + i];
      t1 = t1 + shTan1[base + i];
    }
    let nRaw = cross(t0, t1);
    let nLen = length(nRaw);
    let nrm  = select(vec3<f32>(0.0, 0.0, 0.0), nRaw / nLen, nLen > 0.0);
    outBuf[P.lpOffWords + v * 3u + 0u] = p.x;
    outBuf[P.lpOffWords + v * 3u + 1u] = p.y;
    outBuf[P.lpOffWords + v * 3u + 2u] = p.z;
    outBuf[P.lnOffWords + v * 3u + 0u] = nrm.x;
    outBuf[P.lnOffWords + v * 3u + 1u] = nrm.y;
    outBuf[P.lnOffWords + v * 3u + 2u] = nrm.z;
  }
}
`;

// =====================================================================
// CcGpuLimitCtx -- persistent compute pipeline + static buffers.
//
// Static state (m_vertexLimitInfo stencil tables) is uploaded once at
// ctx.create(device, ccState). Per-frame work is just a srcPositions
// upload + dispatch + readback in ctx.dispatch(srcPositions).
//
// This is the production-shape layout for the rest of Phase 6: each
// compute pass gets its own ctx so the pipeline + read-only storage
// buffers stay resident across frames. Only the per-frame inputs
// (srcPositions + deform deltas) move each frame.
//
// usage:
//   const ctx = CcGpuLimitCtx.create(device, ccState);
//   ...
//   const { positions, normals } = await ctx.dispatch(srcPositions);
//   ...
//   ctx.destroy();  // when the scene unloads
// =====================================================================
class CcGpuLimitCtx {
  // opts.srcBuf -- optional external (shared) srcPositions storage buffer.
  // When provided, this ctx binds it instead of allocating its own (and
  // destroy() leaves it alone -- the caller owns it). Used so reusableCp
  // and limit ctxes can share a single per-frame srcPositions upload
  // instead of writeBufferring the same data twice.
  static create(device, ccState, opts) {
    if (typeof device.createBuffer !== 'function') {
      throw new Error('CcGpuLimitCtx.create: device is not a GPUDevice');
    }
    const li = ccState.patches.m_vertexLimitInfo;
    const numLimitPoints = li.m_limitDataStart.length;
    const numBaseVerts   = ccState.helper0.m_numVerts;

    const dataStartArr     = new Int32Array(li.m_limitDataStart);
    const dataNumArr       = new Int32Array(li.m_limitDataNum);
    const influenceDataArr = new Int32Array(li.m_limitInfluenceData);
    const weightDataPArr   = new Float32Array(li.m_limitWeightDataP);
    const weightDataT0Arr  = new Float32Array(li.m_limitWeightDataT0);
    const weightDataT1Arr  = new Float32Array(li.m_limitWeightDataT1);

    function makeStorageWithData(typedArr, label) {
      const buf = device.createBuffer({
        size: Math.max(typedArr.byteLength, 16),
        usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
        label,
      });
      device.queue.writeBuffer(buf, 0, typedArr.buffer, typedArr.byteOffset, typedArr.byteLength);
      return buf;
    }

    const paramsBuf = device.createBuffer({
      size: 16, usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, label: 'cc_gpu/params',
    });
    device.queue.writeBuffer(paramsBuf, 0, new Uint32Array([numLimitPoints, 0, 0, 0]).buffer);

    // Static stencil tables -- never re-uploaded.
    const startBuf  = makeStorageWithData(dataStartArr,     'cc_gpu/dataStart');
    const numBuf    = makeStorageWithData(dataNumArr,       'cc_gpu/dataNum');
    const inflBuf   = makeStorageWithData(influenceDataArr, 'cc_gpu/influenceData');
    const wpBuf     = makeStorageWithData(weightDataPArr,   'cc_gpu/weightDataP');
    const wt0Buf    = makeStorageWithData(weightDataT0Arr,  'cc_gpu/weightDataT0');
    const wt1Buf    = makeStorageWithData(weightDataT1Arr,  'cc_gpu/weightDataT1');

    // Per-frame srcPositions: persistent (sized at create-time, written
    // each dispatch). When opts.srcBuf is provided, bind that external
    // buffer instead so multiple ctxes share a single upload.
    const externalSrcBuf = opts && opts.srcBuf ? opts.srcBuf : null;
    const srcBuf = externalSrcBuf || device.createBuffer({
      size: Math.max(numBaseVerts * 3 * 4, 16),
      usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
      label: 'cc_gpu/srcPositions',
    });
    const ownsSrcBuf = !externalSrcBuf;

    // Bind group LAYOUT is created here; the bind group itself is built
    // lazily either by bindGpuBlob() (runtime path -- output = gpuBlob,
    // offsets via Params.lp/lnOffWords) or by dispatch() (harness one-shot --
    // output = temp combined buffer).
    const layout = device.createBindGroupLayout({
      entries: [
        { binding: 0, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'uniform'           } },
        { 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: 'read-only-storage' } },
        { binding: 5, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'read-only-storage' } },
        { binding: 6, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'read-only-storage' } },
        { binding: 7, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'read-only-storage' } },
        { binding: 8, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'storage'           } },
      ],
    });
    // A/B-switchable variants for perf experiments. Share the same
    // bind group layout + buffers; only the WGSL module and the per-
    // dispatch (workgroupSize, dispatchCount) shape differ. setVariant()
    // swaps which pipeline encodeCompute() picks.
    const pipelineLayout = device.createPipelineLayout({ bindGroupLayouts: [layout] });
    const variants = {
      'default': {
        pipeline: device.createComputePipeline({
          layout:  pipelineLayout,
          compute: { module: device.createShaderModule({ code: CC_GPU_LIMIT_PN_WGSL }), entryPoint: 'main' },
        }),
        // workgroup_size(64), 1 thread / vert
        dispatchCount: (n) => Math.ceil(n / 64),
      },
      'ws32': {
        pipeline: device.createComputePipeline({
          layout:  pipelineLayout,
          compute: { module: device.createShaderModule({ code: CC_GPU_LIMIT_PN_WS32_WGSL }), entryPoint: 'main' },
        }),
        // workgroup_size(32), 1 thread / vert
        dispatchCount: (n) => Math.ceil(n / 32),
      },
      'coop8': {
        pipeline: device.createComputePipeline({
          layout:  pipelineLayout,
          compute: { module: device.createShaderModule({ code: CC_GPU_LIMIT_PN_COOP_WGSL }), entryPoint: 'main' },
        }),
        // workgroup_size(64), 8 lanes / vert, 8 verts / WG
        dispatchCount: (n) => Math.ceil(n / 8),
      },
    };

    return new CcGpuLimitCtx({
      device, ccState, numLimitPoints, numBaseVerts,
      paramsBuf, startBuf, numBuf, inflBuf, wpBuf, wt0Buf, wt1Buf,
      srcBuf, ownsSrcBuf,
      layout, variants,
      // backwards-compat: the legacy .pipeline field tracks the active
      // variant so any external probe that reaches for it still sees
      // something sane.
      pipeline: variants['default'].pipeline,
      activeVariant: 'default',
      bindGroup: null,  // set by bindGpuBlob() for runtime use
    });
  }

  // Swap the active limit-pn shader variant. Bind group + buffers don't
  // change; only the pipeline reference and the per-dispatch (wgSize,
  // wgCount) shape change. Returns the new variant name.
  setVariant(name) {
    const v = this.variants && this.variants[name];
    if (!v) throw new Error('CcGpuLimitCtx.setVariant: unknown variant ' + name);
    this.activeVariant = name;
    this.pipeline = v.pipeline;
    return name;
  }

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

  // Bind the runtime output target. ccBeginFrameGpu uses topo.gpuBlob;
  // limit's pos/nrm sub-sections are written in-place via the offsets
  // packed into Params. Eliminates the previous outPosBuf/outNrmBuf +
  // 2 copyBufferToBuffer pattern -- and the UAV->COPY_SRC transition
  // barriers those copies needed, which were bottom-of-pipe-attributed
  // to limit_pn's end-of-pass timestamp on Dawn/D3D12.
  bindGpuBlob(outBuf, lpOffWords, lnOffWords) {
    this.device.queue.writeBuffer(this.paramsBuf, 0,
      new Uint32Array([this.numLimitPoints, lpOffWords >>> 0, lnOffWords >>> 0, 0]).buffer);
    this.bindGroup = this.device.createBindGroup({
      layout: this.layout,
      entries: [
        { binding: 0, resource: { buffer: this.paramsBuf } },
        { binding: 1, resource: { buffer: this.startBuf  } },
        { binding: 2, resource: { buffer: this.numBuf    } },
        { binding: 3, resource: { buffer: this.inflBuf   } },
        { binding: 4, resource: { buffer: this.wpBuf     } },
        { binding: 5, resource: { buffer: this.wt0Buf    } },
        { binding: 6, resource: { buffer: this.wt1Buf    } },
        { binding: 7, resource: { buffer: this.srcBuf    } },
        { binding: 8, resource: { buffer: outBuf         } },
      ],
    });
  }

  // Harness one-shot. Writes srcPositions to its srcBuf, dispatches into
  // a temp combined-output buffer (pos at word 0, nrm at word numLimitPoints*3),
  // reads back, splits into { positions, normals }. Self-contained -- doesn't
  // touch the runtime bindGroup state.
  async dispatch(srcPositionsFlat) {
    if (srcPositionsFlat.byteLength > this.numBaseVerts * 3 * 4) {
      throw new Error('CcGpuLimitCtx.dispatch: srcPositions larger than ctx capacity');
    }
    this.device.queue.writeBuffer(this.srcBuf, 0,
      srcPositionsFlat.buffer, srcPositionsFlat.byteOffset, srcPositionsFlat.byteLength);

    const halfWords = this.numLimitPoints * 3;
    const outBytes  = Math.max(halfWords * 2 * 4, 16);
    const tempOutBuf = this.device.createBuffer({
      size: outBytes,
      usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
      label: 'cc_gpu/limitOutOneShot',
    });
    // Params for one-shot: pos at word 0, nrm at word halfWords.
    this.device.queue.writeBuffer(this.paramsBuf, 0,
      new Uint32Array([this.numLimitPoints, 0, halfWords, 0]).buffer);
    const bg = this.device.createBindGroup({
      layout: this.layout,
      entries: [
        { binding: 0, resource: { buffer: this.paramsBuf } },
        { binding: 1, resource: { buffer: this.startBuf  } },
        { binding: 2, resource: { buffer: this.numBuf    } },
        { binding: 3, resource: { buffer: this.inflBuf   } },
        { binding: 4, resource: { buffer: this.wpBuf     } },
        { binding: 5, resource: { buffer: this.wt0Buf    } },
        { binding: 6, resource: { buffer: this.wt1Buf    } },
        { binding: 7, resource: { buffer: this.srcBuf    } },
        { binding: 8, resource: { buffer: tempOutBuf     } },
      ],
    });

    const enc  = this.device.createCommandEncoder({ label: 'cc_gpu/encoder' });
    const pass = enc.beginComputePass({ label: 'cc_gpu/limitPN' });
    pass.setPipeline(this.pipeline);
    pass.setBindGroup(0, bg);
    pass.dispatchWorkgroups(Math.ceil(this.numLimitPoints / 64));
    pass.end();

    const rb = this.device.createBuffer({
      size: outBytes, usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST, label: 'cc_gpu/rbLimit',
    });
    enc.copyBufferToBuffer(tempOutBuf, 0, rb, 0, outBytes);
    this.device.queue.submit([enc.finish()]);

    await rb.mapAsync(GPUMapMode.READ);
    const all = new Float32Array(rb.getMappedRange().slice(0, halfWords * 2 * 4));
    const positions = all.slice(0, halfWords);
    const normals   = all.slice(halfWords, halfWords * 2);
    rb.unmap();
    rb.destroy();
    tempOutBuf.destroy();
    // Reset bindGroup -- if the same ctx is later used via bindGpuBlob,
    // the caller must call bindGpuBlob again. (In practice harness ctxes
    // are short-lived and destroyed right after dispatch.)
    this.bindGroup = null;
    return { positions, normals };
  }

  // Runtime variant. Records the compute pass into the caller's encoder.
  // No writeBuffer, no readback. Caller must have already uploaded
  // srcPositions to the bound srcBuf AND called bindGpuBlob(). limit's
  // outputs land directly in the bound output buffer (e.g. topo.gpuBlob)
  // at the offsets passed to bindGpuBlob -- no separate copy needed.
  encodeCompute(enc, tsOpts) {
    if (!this.bindGroup) throw new Error('CcGpuLimitCtx.encodeCompute: bindGpuBlob() must be called first');
    const beg = { label: 'cc_gpu/limitPN' };
    if (tsOpts && tsOpts.slot >= 0) beg.timestampWrites = { querySet: tsOpts.querySet, beginningOfPassWriteIndex: tsOpts.slot, endOfPassWriteIndex: tsOpts.slot + 1 };
    const pass = enc.beginComputePass(beg);
    pass.setPipeline(this.pipeline);
    pass.setBindGroup(0, this.bindGroup);
    const variant = this.variants[this.activeVariant];
    pass.dispatchWorkgroups(variant.dispatchCount(this.numLimitPoints));
    pass.end();
  }

  destroy() {
    this.paramsBuf.destroy();
    this.startBuf.destroy();  this.numBuf.destroy();
    this.inflBuf.destroy();   this.wpBuf.destroy();
    this.wt0Buf.destroy();    this.wt1Buf.destroy();
    if (this.ownsSrcBuf) this.srcBuf.destroy();
  }
}

// =====================================================================
// gpuComputeCcLimitPN
//
// One-shot wrapper around CcGpuLimitCtx for the harness validation
// gate. Builds the ctx, dispatches once, destroys. Production code
// should use CcGpuLimitCtx directly + dispatch each frame.
//
// `device` is the WebGPU device.
// `ccState` is the cc_render.js state object (carries
// patches.m_vertexLimitInfo + the limit-point count).
// `srcPositionsFlat` is a Float32Array of (numBaseVerts * 3) source
// positions -- AFTER prepMeshForCatmullClark's UV-id duplication
// (same buffer ccBeginFrame consumes).
//
// One-shot: rebuilds the pipeline + bind groups + uploads + dispatches
// + reads back per call. For per-frame use the viewer would persist
// the pipeline + the stencil tables (static) and only re-upload
// srcPositions; this wrapper is for the harness validation gate.
async function gpuComputeCcLimitPN(device, ccState, srcPositionsFlat) {
  const ctx = CcGpuLimitCtx.create(device, ccState);
  try {
    return await ctx.dispatch(srcPositionsFlat);
  } finally {
    ctx.destroy();
  }
}

// =====================================================================
// gpuComputeCcWedgeIter (Phase 6c)
//
// Dispatch the irregular-wedge subdivision compute. One workgroup per
// wedge, runs irregularProcLevel iterations within the workgroup.
// Returns the flat dstWedgePts buffer, comparable to CPU
// state.fullWedgeDstPoints (an array of [x,y,z]).
//
// One-shot for the harness gate; production code would persist
// pipeline + static buffers per CcGpuWedgeCtx (TODO).
// =====================================================================
// =====================================================================
// CC_GPU_REUSABLE_CP_WGSL (Phase 6e)
//
// Per-frame stencil application that builds reusableCpData -- one
// Vec3 per runtime CP. Mirrors ccBeginFrameComputeReusableCpData
// (cc_patcheval.js) on GPU. One workgroup thread per output CP:
//   - For i < numBaseVerts: copy srcPositions[i] directly.
//   - For i >= numBaseVerts: apply the (cpItemStart, cpItemCount,
//     inflData, weightData) stencil to srcPositions.
//
// Foundation for full-GPU ccBeginFrame: the wedge subdiv compute
// currently re-computes reusableCpData on CPU each dispatch + uploads
// it. With this pass, the wedge ctx can chain off the GPU output
// directly (a future refactor; out of scope for the validation gate).
// =====================================================================
const CC_GPU_REUSABLE_CP_WGSL = /* wgsl */ `
// P.cpOffWords -- f32-index offset where outCps[0] should land in
// the bound output storage buffer. Production binds outCps to
// topo.gpuBlob with cpOffWords = vec3DataByteOff/4 + cpDataOffsetV3*3
// so the cp section lands directly in gpuBlob -- eliminates the
// per-frame copyBufferToBuffer(outBuf -> gpuBlob, cpOff) that used
// to follow this pass and whose UAV->COPY_DST barrier on gpuBlob
// was bottom-of-pipe-attributed back to limit_pn's end-of-pass timer.
// Harness one-shot dispatch sets cpOffWords = 0 and binds its own
// temp output buffer.
struct Params {
  numBaseCps    : u32,
  numTotalCps   : u32,
  cpOffWords    : u32,
  _pad0         : u32,
};
@group(0) @binding(0) var<uniform>             P            : Params;
@group(0) @binding(1) var<storage, read>       cpItemStart  : array<i32>;
@group(0) @binding(2) var<storage, read>       cpItemCount  : array<i32>;
@group(0) @binding(3) var<storage, read>       inflData     : array<i32>;
@group(0) @binding(4) var<storage, read>       weightData   : array<f32>;
@group(0) @binding(5) var<storage, read>       srcPositions : array<f32>;
@group(0) @binding(6) var<storage, read_write> outCps       : array<f32>;

@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
  let i = gid.x;
  if (i >= P.numTotalCps) { return; }
  let base = P.cpOffWords + i * 3u;
  if (i < P.numBaseCps) {
    outCps[base + 0u] = srcPositions[i*3u + 0u];
    outCps[base + 1u] = srcPositions[i*3u + 1u];
    outCps[base + 2u] = srcPositions[i*3u + 2u];
    return;
  }
  let cpIndex = i - P.numBaseCps;
  let start   = u32(cpItemStart[cpIndex]);
  let count   = u32(cpItemCount[cpIndex]);
  var px : f32 = 0.0;
  var py : f32 = 0.0;
  var pz : f32 = 0.0;
  for (var k : u32 = 0u; k < count; k = k + 1u) {
    let infl = u32(inflData[start + k]);
    let w    = weightData[start + k];
    px = px + srcPositions[infl*3u + 0u] * w;
    py = py + srcPositions[infl*3u + 1u] * w;
    pz = pz + srcPositions[infl*3u + 2u] * w;
  }
  outCps[base + 0u] = px;
  outCps[base + 1u] = py;
  outCps[base + 2u] = pz;
}
`;

// =====================================================================
// CcGpuReusableCpCtx -- persistent pipeline + static stencil tables.
// Per-frame dispatch writes srcPositions + dispatches + reads back
// reusableCpData. The viewer would chain this with CcGpuWedgeCtx /
// CcGpuLimitCtx in a real per-frame BeginFrame to avoid the
// reusableCpData round-trip through CPU memory.
// =====================================================================
class CcGpuReusableCpCtx {
  // opts.srcBuf -- optional external (shared) srcPositions buffer. See
  // CcGpuLimitCtx.create for rationale.
  static create(device, ccState, opts) {
    const w = ccState.patches.m_weightInflData;
    const numBaseVerts = ccState.helper0.m_numVerts;
    const numTotalCps  = w.m_numBaseCps + w.m_cpItemStart.length;

    function makeStorageWithData(typedArr, label) {
      const buf = device.createBuffer({
        size: Math.max(typedArr.byteLength, 16),
        usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
        label,
      });
      device.queue.writeBuffer(buf, 0, typedArr.buffer, typedArr.byteOffset, typedArr.byteLength);
      return buf;
    }

    const params = device.createBuffer({
      size: 16, usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, label: 'cc_gpu_cp/params',
    });
    // cpOffWords starts at 0; bindGpuBlob() overwrites this when the
    // runtime binds gpuBlob as the output.
    device.queue.writeBuffer(params, 0, new Uint32Array([w.m_numBaseCps, numTotalCps, 0, 0]).buffer);

    const startBuf  = makeStorageWithData(new Int32Array(w.m_cpItemStart),   'cc_gpu_cp/cpItemStart');
    const countBuf  = makeStorageWithData(new Int32Array(w.m_cpItemCount),   'cc_gpu_cp/cpItemCount');
    const inflBuf   = makeStorageWithData(new Int32Array(w.m_inflData),      'cc_gpu_cp/inflData');
    const weightBuf = makeStorageWithData(new Float32Array(w.m_weightData),  'cc_gpu_cp/weightData');

    const externalSrcBuf = opts && opts.srcBuf ? opts.srcBuf : null;
    const srcBuf = externalSrcBuf || device.createBuffer({
      size: Math.max(numBaseVerts * 3 * 4, 16),
      usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
      label: 'cc_gpu_cp/srcPositions',
    });
    const ownsSrcBuf = !externalSrcBuf;
    const outBytes = Math.max(numTotalCps * 3 * 4, 16);

    const layout = device.createBindGroupLayout({
      entries: [
        { binding: 0, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'uniform'           } },
        { 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: 'read-only-storage' } },
        { binding: 5, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'read-only-storage' } },
        { binding: 6, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'storage'           } },
      ],
    });
    const pipeline = device.createComputePipeline({
      layout:  device.createPipelineLayout({ bindGroupLayouts: [layout] }),
      compute: { module: device.createShaderModule({ code: CC_GPU_REUSABLE_CP_WGSL }), entryPoint: 'main' },
    });

    return new CcGpuReusableCpCtx({
      device, numBaseCps: w.m_numBaseCps, numBaseVerts, numTotalCps, outBytes,
      params, startBuf, countBuf, inflBuf, weightBuf,
      srcBuf, ownsSrcBuf,
      layout, pipeline,
      // bindGroup is set lazily by bindGpuBlob() (runtime path) or by
      // dispatch() (harness one-shot). outBuf is null in the runtime
      // path (writes land directly in gpuBlob); dispatch() allocates
      // its own temp output for readback.
      bindGroup: null, outBuf: null,
    });
  }

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

  // Runtime binding: the cp section of topo.gpuBlob is the output
  // target. cpOffWords is an f32-index offset; the shader writes
  // outCps[cpOffWords + i*3 + k]. Eliminates the per-frame copy that
  // used to land an intermediate outBuf into gpuBlob.
  bindGpuBlob(outBuf, cpOffWords) {
    this.device.queue.writeBuffer(this.params, 0,
      new Uint32Array([this.numBaseCps, this.numTotalCps, cpOffWords >>> 0, 0]).buffer);
    this.bindGroup = this.device.createBindGroup({
      layout: this.layout,
      entries: [
        { binding: 0, resource: { buffer: this.params    } },
        { binding: 1, resource: { buffer: this.startBuf  } },
        { binding: 2, resource: { buffer: this.countBuf  } },
        { binding: 3, resource: { buffer: this.inflBuf   } },
        { binding: 4, resource: { buffer: this.weightBuf } },
        { binding: 5, resource: { buffer: this.srcBuf    } },
        { binding: 6, resource: { buffer: outBuf         } },
      ],
    });
  }

  async dispatch(srcPositionsFlat) {
    this.device.queue.writeBuffer(this.srcBuf, 0,
      srcPositionsFlat.buffer, srcPositionsFlat.byteOffset, srcPositionsFlat.byteLength);

    // Harness one-shot: allocate a temp output buffer and bind it at
    // cpOffWords=0 so the shader writes into [0, numTotalCps*3) like
    // it used to.
    const tempOutBuf = this.device.createBuffer({
      size:  this.outBytes,
      usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
      label: 'cc_gpu_cp/outOneShot',
    });
    this.bindGpuBlob(tempOutBuf, 0);

    const enc  = this.device.createCommandEncoder();
    const pass = enc.beginComputePass();
    pass.setPipeline(this.pipeline);
    pass.setBindGroup(0, this.bindGroup);
    pass.dispatchWorkgroups(Math.ceil(this.numTotalCps / 64));
    pass.end();

    const rb = this.device.createBuffer({
      size: this.outBytes, usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST,
    });
    enc.copyBufferToBuffer(tempOutBuf, 0, rb, 0, this.outBytes);
    this.device.queue.submit([enc.finish()]);
    await rb.mapAsync(GPUMapMode.READ);
    const out = new Float32Array(rb.getMappedRange().slice(0, this.numTotalCps * 3 * 4));
    rb.unmap();
    rb.destroy();
    tempOutBuf.destroy();
    // Reset bindGroup -- the caller must re-bind via bindGpuBlob to
    // use this ctx for runtime encodeCompute.
    this.bindGroup = null;
    return out;
  }

  // Runtime variant: encode-only. See CcGpuLimitCtx.encodeCompute.
  encodeCompute(enc, tsOpts) {
    if (!this.bindGroup) throw new Error('CcGpuReusableCpCtx.encodeCompute: bindGpuBlob() must be called first');
    const beg = { label: 'cc_gpu/reusableCp' };
    if (tsOpts && tsOpts.slot >= 0) beg.timestampWrites = { querySet: tsOpts.querySet, beginningOfPassWriteIndex: tsOpts.slot, endOfPassWriteIndex: tsOpts.slot + 1 };
    const pass = enc.beginComputePass(beg);
    pass.setPipeline(this.pipeline);
    pass.setBindGroup(0, this.bindGroup);
    pass.dispatchWorkgroups(Math.ceil(this.numTotalCps / 64));
    pass.end();
  }

  destroy() {
    this.params.destroy();
    this.startBuf.destroy(); this.countBuf.destroy();
    this.inflBuf.destroy();  this.weightBuf.destroy();
    if (this.ownsSrcBuf) this.srcBuf.destroy();
    // outBuf is no longer owned here -- in runtime mode it's gpuBlob
    // (caller-owned); in dispatch() mode it's a temp destroyed in line.
  }
}

// =====================================================================
// CcGpuWedgeCtx -- persistent pipeline + static buffers for the
// irregular-wedge subdivision compute. Same shape as CcGpuLimitCtx.
// Static state (per-wedge config, cpRedirect, irregular edge
// sharpness, per-wedge thread sharpness) is uploaded once at
// create(); each dispatch() writes just reusableCpData + dispatches +
// reads back the per-wedge subdivision output (dstWedgePts).
// =====================================================================
class CcGpuWedgeCtx {
  // Packed multi-wedge port: workgroup_size(64), multiple wedges per
  // workgroup via per-thread (wedgeIdxInGroup, threadCorner) lookup
  // tables that mirror the CPU's IRREGULAR_GROUP_SIZE=64 bin packing.
  // Max valence per wedge = 64.
  //
  // opts.reusableBuf -- optional external storage buffer that already
  // holds reusableCpData (typically CcGpuReusableCpCtx.outBuf). When
  // supplied, wedge binds it directly instead of owning a separate
  // copy, and ccBeginFrameGpu can drop the per-frame
  // copyBufferToBuffer between the reusableCp output and the wedge
  // input -- saves ~numTotalCps*12 bytes/frame plus two state-
  // transition barriers (UAV->COPY_SRC on reusableCp.outBuf,
  // COPY_DST->STORAGE on the previously-owned reusableBuf). The
  // harness one-shot dispatch() path leaves opts unset and gets an
  // owned reusableBuf because it computes reusableCpData on CPU and
  // writeBuffer's it in.
  static create(device, ccState, opts) {
    if (typeof device.createBuffer !== 'function') {
      throw new Error('CcGpuWedgeCtx.create: device is not a GPUDevice');
    }
    const w = ccState.patches.m_weightInflData;
    const numWedges = ccState.patches.m_irregularWedgeInfo.length;
    const numTotalCps = w.m_numBaseCps + w.m_cpItemStart.length;
    const SUBDIV_MAX_SHARPNESS = 10;
    const GROUP_SIZE = 64;
    const MAX_VALENCE = 64;
    const DEAD_LANE = 0xFFFF;

    // Per-wedge config -- same layout as before, indexed by global wedge id.
    const cfgWordsPerWedge = 12;
    const cfgBuf  = new ArrayBuffer(numWedges * cfgWordsPerWedge * 4);
    const cfgU32  = new Uint32Array(cfgBuf);
    const cfgI32  = new Int32Array(cfgBuf);
    const cfgF32  = new Float32Array(cfgBuf);

    const irregLevel = ccState.irregularProcLevel || 4.0;
    const numIterations = Math.max(1, Math.ceil(irregLevel));

    for (let wedge = 0; wedge < numWedges; wedge++) {
      const rii = w.m_runtimeIrregularInfo[wedge];
      if (rii.m_valence > MAX_VALENCE) {
        throw new Error('CcGpuWedgeCtx: wedge ' + wedge + ' valence=' + rii.m_valence + ' exceeds MAX_VALENCE=' + MAX_VALENCE);
      }
      const numPointsPerSlice = (rii.m_isBorder ? 3 : 0) + 1 + rii.m_valence * 12;
      const dstOffset = ccState.wedgeStartPoints[wedge];
      const base = wedge * cfgWordsPerWedge;
      cfgU32[base + 0]  = rii.m_pointStart;
      cfgU32[base + 1]  = rii.m_valence;
      cfgU32[base + 2]  = rii.m_isBorder ? 1 : 0;
      cfgI32[base + 3]  = rii.m_edgeSharpnessStart;
      cfgF32[base + 4]  = rii.m_centerVertexSharpness;
      cfgF32[base + 5]  = rii.m_secondHighestSharpness !== undefined ? rii.m_secondHighestSharpness : rii.m_secondHighestEdgeSharpness;
      cfgF32[base + 6]  = rii.m_thirdHighestSharpness  !== undefined ? rii.m_thirdHighestSharpness  : rii.m_thirdHighestEdgeSharpness;
      cfgU32[base + 7]  = rii.m_centerSemisharpEdge0;
      cfgU32[base + 8]  = rii.m_centerSemisharpEdge1;
      cfgU32[base + 9]  = dstOffset;
      cfgU32[base + 10] = numIterations;
      cfgU32[base + 11] = numPointsPerSlice;
    }

    // Bin-pack wedges into GROUP_SIZE-lane workgroups. Mirrors the CPU
    // bin-packing at cc_topology.js:4685+. Each group covers a contiguous
    // range of global wedge ids; per-thread tables tell each lane which
    // wedge (within the group) and corner it owns. Lanes past the last
    // wedge in a group are marked dead.
    const groupStartArr = [];
    const groupNumArr   = [];
    let curGroupStart = 0;
    let curGroupNum   = 0;
    let lanesUsed     = 0;
    for (let w2 = 0; w2 < numWedges; w2++) {
      const rii = w.m_runtimeIrregularInfo[w2];
      const threadsReq = rii.m_isBorder ? (rii.m_valence + 1) : rii.m_valence;
      if (lanesUsed + threadsReq > GROUP_SIZE) {
        groupStartArr.push(curGroupStart);
        groupNumArr.push(curGroupNum);
        curGroupStart = w2;
        curGroupNum   = 0;
        lanesUsed     = 0;
      }
      lanesUsed += threadsReq;
      curGroupNum++;
    }
    if (curGroupNum > 0) {
      groupStartArr.push(curGroupStart);
      groupNumArr.push(curGroupNum);
    }
    const numGroups = groupStartArr.length;
    const groupStart = new Uint32Array(groupStartArr);
    // Per-thread lookup: high 16 bits = wedgeIdxInGroup (DEAD_LANE = dead),
    // low 16 bits = threadCorner.
    const perThread = new Uint32Array(numGroups * GROUP_SIZE);
    const perThreadEdgeShp = new Float32Array(numGroups * GROUP_SIZE);
    for (let g = 0; g < numGroups; g++) {
      let lane = 0;
      for (let wig = 0; wig < groupNumArr[g]; wig++) {
        const wedge = groupStartArr[g] + wig;
        const rii = w.m_runtimeIrregularInfo[wedge];
        const numThreadsForWedge = rii.m_isBorder ? (rii.m_valence + 1) : rii.m_valence;
        for (let tc = 0; tc < numThreadsForWedge; tc++) {
          perThread[g * GROUP_SIZE + lane] = (wig << 16) | (tc & 0xFFFF);
          // Per-thread edge sharpness for this lane. For borders, the
          // phantom lane (tc == valence) has no sharpness entry --
          // leave 0.
          if (rii.m_edgeSharpnessStart >= 0 && tc < rii.m_valence) {
            perThreadEdgeShp[g * GROUP_SIZE + lane] =
              w.m_irregularEdgeSharpness[rii.m_edgeSharpnessStart + tc];
          }
          lane++;
        }
      }
      for (; lane < GROUP_SIZE; lane++) {
        perThread[g * GROUP_SIZE + lane] = (DEAD_LANE << 16);
      }
    }

    function makeStorageWithData(typedArr, label) {
      const buf = device.createBuffer({
        size: Math.max(typedArr.byteLength, 16),
        usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
        label,
      });
      device.queue.writeBuffer(buf, 0, typedArr.buffer, typedArr.byteOffset, typedArr.byteLength);
      return buf;
    }

    const globalU = device.createBuffer({
      size: 32, usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, label: 'cc_gpu_wedge/U',
    });
    // cpOffWords / wedgeOffWords (slots 5 + 6) start at 0; bindGpuBlob
    // overwrites them when the runtime binds gpuBlob as the storage
    // target.
    device.queue.writeBuffer(globalU, 0,
      new Uint32Array([numWedges, w.m_numRegularPoints || 0, 0, numGroups, numIterations, 0, 0, 0]).buffer);
    device.queue.writeBuffer(globalU, 8, new Float32Array([SUBDIV_MAX_SHARPNESS]).buffer);

    // Static buffers (uploaded once).
    const cfgGpuBuf         = makeStorageWithData(new Uint8Array(cfgBuf), 'cc_gpu_wedge/cfg');
    const regularAndIrreg   = makeStorageWithData(new Int32Array(w.m_regularAndIrregularPoints), 'cc_gpu_wedge/regAndIrreg');
    const perThreadEdgeBuf  = makeStorageWithData(perThreadEdgeShp, 'cc_gpu_wedge/perThreadEdgeShp');
    const groupStartBuf     = makeStorageWithData(groupStart, 'cc_gpu_wedge/groupStart');
    const perThreadBuf      = makeStorageWithData(perThread, 'cc_gpu_wedge/perThread');

    // Output buffer sizing (still tracked even though we don't own
    // one): totalDstPts is the cumulative point count across all
    // wedges at MAX_IRREG_LEVEL preallocation. The runtime path lands
    // those points inside topo.gpuBlob's vec3 section; the harness
    // dispatch() path allocates a temp combined buffer big enough to
    // hold cp data + wedge slice data side by side.
    let totalDstPts = 0;
    for (let i = 0; i < numWedges; i++) totalDstPts += ccState.wedgeNumDstPoints[i];
    const dstBytes = Math.max(totalDstPts * 3 * 4, 16);

    const layout = device.createBindGroupLayout({
      entries: [
        { binding: 0, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'uniform'           } },
        { binding: 1, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'read-only-storage' } },
        { binding: 2, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'read-only-storage' } },
        // binding 3 (reusableCpData) was dropped; cp data is now read
        // from binding 7 with U.cpOffWords offset. WebGPU rejects
        // aliasing a buffer as read-only-storage + storage in the
        // same dispatch, hence the single-binding scheme.
        { binding: 4, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'read-only-storage' } },
        { binding: 5, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'read-only-storage' } },
        { binding: 6, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'read-only-storage' } },
        { binding: 7, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'storage'           } },
      ],
    });
    const pipeline = device.createComputePipeline({
      layout:  device.createPipelineLayout({ bindGroupLayouts: [layout] }),
      compute: { module: device.createShaderModule({ code: CC_GPU_WEDGE_ITER_WGSL }), entryPoint: 'main' },
    });

    return new CcGpuWedgeCtx({
      device, ccState, numWedges, numGroups, numTotalCps, totalDstPts, dstBytes,
      globalU, cfgGpuBuf, regularAndIrreg, perThreadEdgeBuf,
      groupStartBuf, perThreadBuf,
      layout, pipeline,
      // bindGroup is set lazily by bindGpuBlob() (runtime) or by
      // dispatch() (harness one-shot). dstBuf is null in the runtime
      // path -- writes land directly in gpuBlob.
      bindGroup: null, dstBuf: null,
      _weightInflData: w,
    });
  }

  // Runtime binding: cp data lives in outBuf at cpOffWords (= where
  // CcGpuReusableCpCtx writes), wedge slice data lands in outBuf at
  // wedgeOffWords. Production always sets outBuf = topo.gpuBlob.
  bindGpuBlob(outBuf, cpOffWords, wedgeOffWords) {
    this.device.queue.writeBuffer(this.globalU, 20,
      new Uint32Array([cpOffWords >>> 0, wedgeOffWords >>> 0]).buffer);
    this.bindGroup = this.device.createBindGroup({
      layout: this.layout,
      entries: [
        { binding: 0, resource: { buffer: this.globalU            } },
        { binding: 1, resource: { buffer: this.cfgGpuBuf          } },
        { binding: 2, resource: { buffer: this.regularAndIrreg    } },
        // binding 3 intentionally skipped (see layout comment).
        { binding: 4, resource: { buffer: this.perThreadEdgeBuf   } },
        { binding: 5, resource: { buffer: this.groupStartBuf      } },
        { binding: 6, resource: { buffer: this.perThreadBuf       } },
        { binding: 7, resource: { buffer: outBuf                  } },
      ],
    });
  }

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

  // Per-frame: compute reusableCpData from srcPositions (mirrors
  // ccBeginFrameComputeReusableCpData), upload, dispatch, read back
  // dstWedgePts. Returns a Float32Array of flat-xyz.
  async dispatch(srcPositionsFlat) {
    const w = this._weightInflData;
    const reusableFlat = new Float32Array(this.numTotalCps * 3);
    for (let i = 0; i < w.m_numBaseCps; i++) {
      reusableFlat[i*3+0] = srcPositionsFlat[i*3+0];
      reusableFlat[i*3+1] = srcPositionsFlat[i*3+1];
      reusableFlat[i*3+2] = srcPositionsFlat[i*3+2];
    }
    for (let i = 0; i < w.m_cpItemStart.length; i++) {
      const dst = w.m_numBaseCps + i;
      const start = w.m_cpItemStart[i];
      const count = w.m_cpItemCount[i];
      let x = 0, y = 0, z = 0;
      for (let k = 0; k < count; k++) {
        const sIdx = w.m_inflData[start + k];
        const wt   = w.m_weightData[start + k];
        x += srcPositionsFlat[sIdx*3+0] * wt;
        y += srcPositionsFlat[sIdx*3+1] * wt;
        z += srcPositionsFlat[sIdx*3+2] * wt;
      }
      reusableFlat[dst*3+0] = x; reusableFlat[dst*3+1] = y; reusableFlat[dst*3+2] = z;
    }
    // Harness one-shot: pack cp data + wedge slice region into a
    // single temp buffer so the dispatch sees the same gpuBlob-style
    // layout the runtime uses. cp data occupies the first
    // numTotalCps * 12 bytes; wedge slice data starts right after at
    // wedgeOffWords = numTotalCps * 3.
    const cpBytes  = this.numTotalCps * 3 * 4;
    const tempSize = Math.max(cpBytes + this.dstBytes, 16);
    const tempBuf = this.device.createBuffer({
      size:  tempSize,
      usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST,
      label: 'cc_gpu_wedge/tempCombined',
    });
    this.device.queue.writeBuffer(tempBuf, 0,
      reusableFlat.buffer, reusableFlat.byteOffset, reusableFlat.byteLength);
    this.bindGpuBlob(tempBuf, 0, this.numTotalCps * 3);

    const enc  = this.device.createCommandEncoder();
    const pass = enc.beginComputePass();
    pass.setPipeline(this.pipeline);
    pass.setBindGroup(0, this.bindGroup);
    pass.dispatchWorkgroups(this.numGroups);
    pass.end();

    const rb = this.device.createBuffer({
      size: this.dstBytes, usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST,
    });
    enc.copyBufferToBuffer(tempBuf, cpBytes, rb, 0, this.dstBytes);
    this.device.queue.submit([enc.finish()]);
    await rb.mapAsync(GPUMapMode.READ);
    const out = new Float32Array(rb.getMappedRange().slice(0, this.totalDstPts * 3 * 4));
    rb.unmap();
    rb.destroy();
    tempBuf.destroy();
    // Reset bindGroup -- the caller must rebind via bindGpuBlob to
    // use this ctx for runtime encodeCompute.
    this.bindGroup = null;
    return out;
  }

  // Runtime variant: encode-only. The bind group is set up by
  // bindGpuBlob() pointing at topo.gpuBlob with the cp-data and
  // wedge-data offsets baked into U. Reading cp data and writing
  // wedge slice data both happen through binding 7 (the same
  // gpuBlob) -- they target non-overlapping regions so there's no
  // hazard within the dispatch.
  encodeCompute(enc, tsOpts) {
    if (!this.bindGroup) throw new Error('CcGpuWedgeCtx.encodeCompute: bindGpuBlob() must be called first');
    const beg = { label: 'cc_gpu/wedge' };
    if (tsOpts && tsOpts.slot >= 0) beg.timestampWrites = { querySet: tsOpts.querySet, beginningOfPassWriteIndex: tsOpts.slot, endOfPassWriteIndex: tsOpts.slot + 1 };
    const pass = enc.beginComputePass(beg);
    pass.setPipeline(this.pipeline);
    pass.setBindGroup(0, this.bindGroup);
    pass.dispatchWorkgroups(this.numGroups);
    pass.end();
  }

  // Live update for the runtime numIterations uniform. The shader's
  // for-loop bound is U.numIterations, so this controls how many CC
  // refinement passes get computed per wedge per frame. Dst buffer is
  // pre-sized for MAX_IRREG_LEVEL slices, so any 1..MAX_IRREG_LEVEL
  // value is in-range without reallocation. Called by the irreg-level
  // slider input handler.
  //
  // GlobalU layout (cc_gpu.js:906):
  //   [0] numWedges (u32)
  //   [1] numRegularPoints (u32)
  //   [2] subdivMaxSharpness (f32)
  //   [3] numGroups (u32)
  //   [4] numIterations (u32)   <-- byte 16
  setNumIterations(n) {
    this.device.queue.writeBuffer(this.globalU, 16, new Uint32Array([n >>> 0]).buffer);
  }
  destroy() {
    this.globalU.destroy(); this.cfgGpuBuf.destroy(); this.regularAndIrreg.destroy();
    this.perThreadEdgeBuf.destroy();
    this.groupStartBuf.destroy(); this.perThreadBuf.destroy();
    // reusableBuf + dstBuf no longer owned -- runtime mode points at
    // gpuBlob (caller-owned), harness dispatch() destroys its own
    // temp buffer in line.
  }
}

// One-shot wrapper around CcGpuWedgeCtx for the harness validation
// gate (kept for backward-compat with the existing __viewerTest hook).
async function gpuComputeCcWedgeIter(device, ccState, srcPositionsFlat) {
  const ctx = CcGpuWedgeCtx.create(device, ccState);
  try {
    return await ctx.dispatch(srcPositionsFlat);
  } finally {
    ctx.destroy();
  }
}

// Legacy alias for callers that only want positions back.
async function gpuComputeCcLimitPositions(device, ccState, srcPositionsFlat) {
  const r = await gpuComputeCcLimitPN(device, ccState, srcPositionsFlat);
  return r.positions;
}

// =====================================================================
// CC_GPU_WEDGE_ITER_WGSL  (packed-64 multi-wedge port)
//
// Per-frame irregular-wedge subdivision compute. Workgroup_size = 64;
// multiple wedges packed per workgroup via per-thread (wedgeIdxInGroup,
// threadCorner) lookup, mirroring amolite's CPU IRREGULAR_GROUP_SIZE=64
// bin-packing (cc_topology.js irregularWedgeRunIteration). Max valence
// per wedge = 64 (a single valence-64 wedge fills the whole workgroup).
//
// Bindings (8 storage total = WebGPU min-spec cap):
//   0 (uniform) U: GlobalU
//   1 (storage read) perWedgeCfg[]
//   2 (storage read) regularAndIrreg[]
//   3 (storage read) reusableCpData[]
//   4 (storage read) perThreadEdgeShp[]   -- f32 per thread, sized numGroups * 64
//   5 (storage read) groupStart[]         -- u32 per group, global-wedge index of first wedge
//   6 (storage read) perThread[]          -- packed u32 per thread: high 16=wedgeIdxInGroup
//                                            (0xFFFFu = dead lane), low 16=threadCorner
//   7 (storage rw)  dstWedgePts[]
//
// Workgroup shared memory layout (= 14,592 bytes; min-spec cap 16,384):
//   innerPrev/innerCurr/outer : array<f32, 64 * 6 * 3>   = 4,608 bytes each
//   centerPrev/centerCurr     : array<f32, MAX_WEDGES_PER_GROUP * 3>  (MAX = 32)
//
// Lane addressing:
//   T   = lid.x in [0, 64)                       -- absolute lane in WG
//   wig = wedgeIdxInGroup in [0, 32)
//   tc  = threadCorner in [0, valence + isBorder)
//   firstThreadForWedge (FTFW) = T - tc          -- lane where THIS wedge starts
//   slotOf(lane, k) = 6 * lane + k               -- inner/outer slot for lane
//   centerSlot(wig) = 3 * wig                    -- center slot for wedge
// =====================================================================
const CC_GPU_WEDGE_ITER_WGSL = /* wgsl */ `
struct WedgeCfg {
  pointStart                : u32,
  valence                   : u32,
  isBorder                  : u32,
  edgeSharpnessStart        : i32,
  centerVertexSharpness     : f32,
  secondHighestSharpness    : f32,
  thirdHighestSharpness     : f32,
  centerSemisharpEdge0      : u32,
  centerSemisharpEdge1      : u32,
  dstOffset                 : u32,
  numIterations             : u32,
  numPointsPerSlice         : u32,
};

// cpOffWords / wedgeOffWords are f32-index offsets into dstWedgePts
// (= topo.gpuBlob in the runtime path). The runtime calls
// bindGpuBlob(buf, cpOffWords, wedgeOffWords) which writes both fields
// into U. Eliminates the two copyBufferToBuffer calls (reusableCp ->
// gpuBlob, wedge -> gpuBlob) that used to follow this pass and whose
// UAV -> COPY_DST barriers on gpuBlob were bottom-of-pipe-attributed
// back to limit_pn's end-of-pass timer.
struct GlobalU {
  numWedges          : u32,
  numRegularPoints   : u32,
  subdivMaxSharpness : f32,
  numGroups          : u32,
  numIterations      : u32,
  cpOffWords         : u32,
  wedgeOffWords      : u32,
  _pad0              : u32,
};

// Binding 3 (reusableCpData) was dropped: the wedge ctx now reads cp
// data from dstWedgePts (= gpuBlob) at U.cpOffWords. WebGPU rejects
// aliasing a buffer as read-only-storage + storage in the same
// dispatch, so cp reads and wedge writes share the one read_write
// binding -- the two regions don't overlap, so there's no hazard.
@group(0) @binding(0)  var<uniform>                                    U                : GlobalU;
@group(0) @binding(1)  var<storage, read>                              perWedgeCfg      : array<WedgeCfg>;
@group(0) @binding(2)  var<storage, read>                              regularAndIrreg  : array<i32>;
@group(0) @binding(4)  var<storage, read>                              perThreadEdgeShp : array<f32>;
@group(0) @binding(5)  var<storage, read>                              groupStart       : array<u32>;
@group(0) @binding(6)  var<storage, read>                              perThread        : array<u32>;
@group(0) @binding(7)  var<storage, read_write>                        dstWedgePts      : array<f32>;

const MAX_WEDGES_PER_GROUP : u32 = 32u;
const DEAD_LANE            : u32 = 0xFFFFu;

var<workgroup> innerPrev  : array<f32, 1152>;
var<workgroup> innerCurr  : array<f32, 1152>;
var<workgroup> outer      : array<f32, 1152>;
var<workgroup> centerPrev : array<f32, 96>;
var<workgroup> centerCurr : array<f32, 96>;

fn fetchSrc(off: u32) -> vec3<f32> {
  let src = u32(regularAndIrreg[U.numRegularPoints + off]);
  let base = U.cpOffWords + src * 3u;
  return vec3<f32>(
    dstWedgePts[base + 0u],
    dstWedgePts[base + 1u],
    dstWedgePts[base + 2u]);
}

fn loadInner(buf: u32, slot: u32) -> vec3<f32> {
  if (buf == 0u) {
    return vec3<f32>(innerPrev[slot * 3u + 0u], innerPrev[slot * 3u + 1u], innerPrev[slot * 3u + 2u]);
  }
  return vec3<f32>(innerCurr[slot * 3u + 0u], innerCurr[slot * 3u + 1u], innerCurr[slot * 3u + 2u]);
}
fn storeInner(buf: u32, slot: u32, v: vec3<f32>) {
  if (buf == 0u) {
    innerPrev[slot * 3u + 0u] = v.x; innerPrev[slot * 3u + 1u] = v.y; innerPrev[slot * 3u + 2u] = v.z;
  } else {
    innerCurr[slot * 3u + 0u] = v.x; innerCurr[slot * 3u + 1u] = v.y; innerCurr[slot * 3u + 2u] = v.z;
  }
}
fn loadOuter(slot: u32) -> vec3<f32> {
  return vec3<f32>(outer[slot * 3u + 0u], outer[slot * 3u + 1u], outer[slot * 3u + 2u]);
}
fn storeOuter(slot: u32, v: vec3<f32>) {
  outer[slot * 3u + 0u] = v.x; outer[slot * 3u + 1u] = v.y; outer[slot * 3u + 2u] = v.z;
}
fn loadCenter(buf: u32, wig: u32) -> vec3<f32> {
  let s = 3u * wig;
  if (buf == 0u) { return vec3<f32>(centerPrev[s + 0u], centerPrev[s + 1u], centerPrev[s + 2u]); }
  return vec3<f32>(centerCurr[s + 0u], centerCurr[s + 1u], centerCurr[s + 2u]);
}
fn storeCenter(buf: u32, wig: u32, v: vec3<f32>) {
  let s = 3u * wig;
  if (buf == 0u) { centerPrev[s + 0u] = v.x; centerPrev[s + 1u] = v.y; centerPrev[s + 2u] = v.z; }
  else           { centerCurr[s + 0u] = v.x; centerCurr[s + 1u] = v.y; centerCurr[s + 2u] = v.z; }
}

fn saturateF(x: f32) -> f32 { return clamp(x, 0.0, 1.0); }

// lane T's 6 inner/outer slots: slot = 6 * lane + k
fn slotOf(lane: u32, k: u32) -> u32 { return 6u * lane + k; }

@compute @workgroup_size(64)
fn main(@builtin(workgroup_id) wid : vec3<u32>,
        @builtin(local_invocation_id) lid : vec3<u32>) {
  let groupId = wid.x;
  if (groupId >= U.numGroups) { return; }
  let T = lid.x;

  // Per-thread lookup: pack = (wedgeIdxInGroup << 16) | threadCorner;
  // wig == DEAD_LANE means this lane has no wedge in this group.
  let pack = perThread[groupId * 64u + T];
  let wig  = pack >> 16u;
  let threadCorner = pack & 0xFFFFu;
  let isDead = (wig == DEAD_LANE);

  // ALL threads sit through the iteration so workgroupBarrier is uniform-
  // control-flow. Dead lanes simply don't read/write.
  var cfg : WedgeCfg;
  var valence : u32 = 0u;
  var isBorder : bool = false;
  var numThreadsForWedge : u32 = 0u;
  var isLive : bool = false;
  var isPhantom : bool = false;
  var firstThreadForWedge : u32 = 0u;
  var currCorner : u32 = 0u;
  var nextCorner : u32 = 0u;
  var prevCorner : u32 = 0u;
  if (!isDead) {
    let wedge = groupStart[groupId] + wig;
    cfg      = perWedgeCfg[wedge];
    valence  = cfg.valence;
    isBorder = cfg.isBorder == 1u;
    numThreadsForWedge = select(valence, valence + 1u, isBorder);
    isLive    = true;
    isPhantom = isBorder && (threadCorner == valence);
    firstThreadForWedge = T - threadCorner;
    currCorner = firstThreadForWedge + threadCorner;
    nextCorner = firstThreadForWedge + ((threadCorner + 1u) % numThreadsForWedge);
    prevCorner = firstThreadForWedge + ((threadCorner + numThreadsForWedge - 1u) % numThreadsForWedge);
  }

  // ---- Init ----
  if (isLive) {
    let baseOffset = cfg.pointStart + 1u + 6u * threadCorner;
    if (isPhantom) {
      let v0 = fetchSrc(baseOffset + 0u);
      let v1 = fetchSrc(baseOffset + 1u);
      storeInner(0u, slotOf(T, 0u), v0);
      storeInner(0u, slotOf(T, 1u), vec3<f32>(-13.0));
      storeInner(0u, slotOf(T, 2u), vec3<f32>(-14.0));
      storeInner(0u, slotOf(T, 3u), v1);
      storeInner(0u, slotOf(T, 4u), vec3<f32>(-23.0));
      storeInner(0u, slotOf(T, 5u), vec3<f32>(-24.0));
    } else {
      for (var k: u32 = 0u; k < 6u; k = k + 1u) {
        storeInner(0u, slotOf(T, k), fetchSrc(baseOffset + k));
      }
    }
    if (threadCorner == 0u) {
      storeCenter(0u, wig, fetchSrc(cfg.pointStart + 0u));
    }
  }
  workgroupBarrier();

  // ---- Iteration loop ----
  // amolite ping-pongs: prevBuf = 0 on entry; after each iter swap.
  var prevBuf : u32 = 0u;
  var currBuf : u32 = 1u;

  // Loop bound from a uniform so workgroupBarrier sits in uniform
  // control flow (cfg.numIterations is per-wedge in the buffer but
  // identical across wedges -- WGSL can't see that without a uniform).
  for (var iter : u32 = 0u; iter < U.numIterations; iter = iter + 1u) {
    // ---- Reset target buffers to debugging sentinels (each LIVE lane
    //      clears its own 6 slots; dead lanes skip) ----
    if (isLive) {
      for (var k: u32 = 0u; k < 6u; k = k + 1u) {
        storeInner(currBuf, slotOf(T, k), vec3<f32>(-2000.0));
        storeOuter(slotOf(T, k), vec3<f32>(-3000.0));
      }
      if (threadCorner == 0u) { storeCenter(currBuf, wig, vec3<f32>(-1000.0)); }
    }
    workgroupBarrier();

    // currCorner/nextCorner/prevCorner are pre-computed LANE indices
    // (firstThreadForWedge + within-wedge-corner). slotOf() then maps
    // each lane to its 6-slot block in inner/outer shared mem.

    let edgeBaseSharpness = select(0.0, perThreadEdgeShp[groupId * 64u + T], isLive);
    let edgeCurrSharpnessT = saturateF(select(edgeBaseSharpness - f32(iter), 10.0, edgeBaseSharpness >= 10.0));
    let isEdgeBorder = isBorder && (threadCorner == 0u || threadCorner == valence);

    // ---- Pass 1: face point f11 ----
    if (isLive) {
      let p00 = loadCenter(prevBuf, wig);
      let p02 = loadInner(prevBuf, slotOf(nextCorner, 0u));
      let p04 = loadInner(prevBuf, slotOf(nextCorner, 3u));
      let p20 = loadInner(prevBuf, slotOf(currCorner, 0u));
      let p22 = loadInner(prevBuf, slotOf(currCorner, 1u));
      let p24 = loadInner(prevBuf, slotOf(currCorner, 2u));
      let p40 = loadInner(prevBuf, slotOf(currCorner, 3u));
      let p42 = loadInner(prevBuf, slotOf(currCorner, 4u));
      let p44 = loadInner(prevBuf, slotOf(currCorner, 5u));

      let f11 = 0.25 * (p00 + p02 + p20 + p22);
      let f13 = 0.25 * (p02 + p04 + p22 + p24);
      let f31 = 0.25 * (p20 + p22 + p40 + p42);
      let f33 = 0.25 * (p22 + p24 + p42 + p44);

      storeInner(currBuf, slotOf(currCorner, 1u), f11);
      storeOuter(slotOf(currCorner, 5u), f13);
      storeOuter(slotOf(currCorner, 1u), f31);
      storeOuter(slotOf(currCorner, 3u), f33);
    }
    workgroupBarrier();

    // ---- Pass 2: edge points e10/e12/e21/e23/e30/e32 ----
    if (isLive) {
      let p00 = loadCenter(prevBuf, wig);
      let p02 = loadInner(prevBuf, slotOf(nextCorner, 0u));
      let p04 = loadInner(prevBuf, slotOf(nextCorner, 3u));
      let p20 = loadInner(prevBuf, slotOf(currCorner, 0u));
      let p22 = loadInner(prevBuf, slotOf(currCorner, 1u));
      let p24 = loadInner(prevBuf, slotOf(currCorner, 2u));
      let p40 = loadInner(prevBuf, slotOf(currCorner, 3u));
      let p42 = loadInner(prevBuf, slotOf(currCorner, 4u));
      let p44 = loadInner(prevBuf, slotOf(currCorner, 5u));
      let f11 = loadInner(currBuf, slotOf(currCorner, 1u));
      let f13 = loadOuter(slotOf(currCorner, 5u));
      let f31 = loadOuter(slotOf(currCorner, 1u));
      let f33 = loadOuter(slotOf(currCorner, 3u));

      let e10_sharp = 0.5 * (p00 + p20);
      let e30_sharp = 0.5 * (p20 + p40);
      var e10 : vec3<f32>;
      var e30 : vec3<f32>;
      if (isEdgeBorder) {
        e10 = e10_sharp;
        e30 = e30_sharp;
      } else {
        let f1_n1 = loadInner(currBuf, slotOf(prevCorner, 1u));
        let f3_n1 = loadOuter(slotOf(prevCorner, 5u));
        let e10_smooth = 0.25 * (p00 + p20 + f1_n1 + f11);
        let e30_smooth = 0.25 * (p20 + p40 + f3_n1 + f31);
        e10 = mix(e10_smooth, e10_sharp, edgeCurrSharpnessT);
        e30 = mix(e30_smooth, e30_sharp, edgeCurrSharpnessT);
      }
      let e12 = 0.25 * (f11 + f13 + p02 + p22);
      let e32 = 0.25 * (f31 + f33 + p22 + p42);
      let e21 = 0.25 * (p20 + p22 + f11 + f31);
      let e23 = 0.25 * (p22 + p24 + f13 + f33);

      storeInner(currBuf, slotOf(currCorner, 0u), e10);
      storeInner(currBuf, slotOf(currCorner, 2u), e12);
      storeInner(currBuf, slotOf(currCorner, 4u), e21);
      storeOuter(slotOf(currCorner, 4u), e23);
      storeOuter(slotOf(currCorner, 0u), e30);
      storeOuter(slotOf(currCorner, 2u), e32);
    }
    workgroupBarrier();

    // ---- Pass 3: vertex points v20/v22 ----
    if (isLive) {
      let p00   = loadCenter(prevBuf, wig);
      let p02   = loadInner(prevBuf, slotOf(nextCorner, 0u));
      let p2_n2 = loadInner(prevBuf, slotOf(prevCorner, 1u));
      let p20   = loadInner(prevBuf, slotOf(currCorner, 0u));
      let p22   = loadInner(prevBuf, slotOf(currCorner, 1u));
      let p24   = loadInner(prevBuf, slotOf(currCorner, 2u));
      let p40   = loadInner(prevBuf, slotOf(currCorner, 3u));
      let p42   = loadInner(prevBuf, slotOf(currCorner, 4u));
      let f11   = loadInner(currBuf, slotOf(currCorner, 1u));
      let f13   = loadOuter(slotOf(currCorner, 5u));
      let f31   = loadOuter(slotOf(currCorner, 1u));
      let f33   = loadOuter(slotOf(currCorner, 3u));

      let inv16 = 0.0625;
      // amolite quirk: v20_sharp uses (p00, p40, p20) -- previous edge vertices, not e10/e30.
      let v20_sharp = 0.125 * (p00 + 6.0 * p20 + p40);
      var v20 : vec3<f32>;
      if (isEdgeBorder) {
        v20 = v20_sharp;
      } else {
        let f1_n1 = loadInner(currBuf, slotOf(prevCorner, 1u));
        let f3_n1 = loadOuter(slotOf(prevCorner, 5u));
        let v20_smooth = 0.5 * p20 +
                         inv16 * (p00 + p40 + p2_n2 + p22) +
                         inv16 * (f1_n1 + f3_n1 + f11 + f31);
        v20 = mix(v20_smooth, v20_sharp, edgeCurrSharpnessT);
      }
      let v22 = 0.5 * p22 +
                inv16 * (p20 + p24 + p02 + p42) +
                inv16 * (f11 + f31 + f13 + f33);

      storeInner(currBuf, slotOf(currCorner, 3u), v20);
      storeInner(currBuf, slotOf(currCorner, 5u), v22);
    }
    workgroupBarrier();

    // ---- Pass 4: center (per-wedge first thread) ----
    if (isLive && threadCorner == 0u) {
      let centerVal = loadCenter(prevBuf, wig);
      // GetCenterVertexEdgeSharpnessInfoHelper (per-frame variant).
      let s2nd  = cfg.secondHighestSharpness;
      let s3rd  = cfg.thirdHighestSharpness;
      let sCtr  = cfg.centerVertexSharpness;
      let secondInfl = saturateF(select(s2nd - f32(iter), 1.0, s2nd >= U.subdivMaxSharpness));
      let thirdInfl  = saturateF(select(s3rd - f32(iter), 1.0, s3rd >= U.subdivMaxSharpness));
      let vertexInfl = saturateF(select(sCtr - f32(iter), 1.0, sCtr >= U.subdivMaxSharpness));
      let cornerInfl = max(thirdInfl, vertexInfl);
      let edgeInfl   = secondInfl;
      let wCorner = cornerInfl;
      let wEdge   = (1.0 - cornerInfl) * edgeInfl;
      let wSmooth = (1.0 - cornerInfl) * (1.0 - edgeInfl);

      // sumEdge/sumFace walk the wedge's corners; the inner-buf slot
      // for corner c lives at LANE (firstThreadForWedge + c).
      var sumEdge = vec3<f32>(0.0);
      var sumFace = vec3<f32>(0.0);
      for (var c : u32 = 0u; c < valence; c = c + 1u) {
        sumEdge = sumEdge + loadInner(prevBuf, slotOf(firstThreadForWedge + c, 0u));
        sumFace = sumFace + loadInner(currBuf, slotOf(firstThreadForWedge + c, 1u));
      }
      let scaleC  = (f32(valence) - 2.0) / f32(valence);
      let scaleEF = 1.0 / (f32(valence) * f32(valence));
      let dst_smooth = scaleC * centerVal + scaleEF * sumEdge + scaleEF * sumFace;

      let valEdge0 = loadInner(prevBuf, slotOf(firstThreadForWedge + cfg.centerSemisharpEdge0, 0u));
      let valEdge1 = loadInner(prevBuf, slotOf(firstThreadForWedge + cfg.centerSemisharpEdge1, 0u));
      let dst_edge = 0.125 * (valEdge0 + 6.0 * centerVal + valEdge1);
      let dst_corner = centerVal;

      let dst = wCorner * dst_corner + wEdge * dst_edge + wSmooth * dst_smooth;
      storeCenter(currBuf, wig, dst);
    }
    workgroupBarrier();

    // ---- SavePoints (iter writes to slice iter) ----
    if (isLive) {
      let v00 = loadInner(currBuf, slotOf(currCorner, 0u));
      let v01 = loadInner(currBuf, slotOf(currCorner, 1u));
      let v02 = loadInner(currBuf, slotOf(currCorner, 2u));
      let v03 = loadOuter(slotOf(currCorner, 5u));
      let v10 = loadInner(currBuf, slotOf(currCorner, 3u));
      let v11 = loadInner(currBuf, slotOf(currCorner, 4u));
      let v12 = loadInner(currBuf, slotOf(currCorner, 5u));
      let v13 = loadOuter(slotOf(currCorner, 4u));
      let v20 = loadOuter(slotOf(currCorner, 0u));
      let v21 = loadOuter(slotOf(currCorner, 1u));
      let v22 = loadOuter(slotOf(currCorner, 2u));
      let v23 = loadOuter(slotOf(currCorner, 3u));

      let dstFullOffset = cfg.dstOffset + cfg.numPointsPerSlice * iter;
      // dstFullOffset is in points (vec3 units); add U.wedgeOffWords
      // (f32 index) after the * 3u to convert to the final f32 index
      // into dstWedgePts (= topo.gpuBlob).
      if (threadCorner == 0u) {
        let c = loadCenter(currBuf, wig);
        let off3 = U.wedgeOffWords + (dstFullOffset + 0u) * 3u;
        dstWedgePts[off3 + 0u] = c.x;
        dstWedgePts[off3 + 1u] = c.y;
        dstWedgePts[off3 + 2u] = c.z;
      }
      // 12 packed Vec3s per corner -- 3 packed for phantom border lane.
      // Offset uses threadCorner (within-wedge), not the LANE index.
      if (isPhantom) {
        let off = dstFullOffset + 1u + 12u * threadCorner;
        let base = U.wedgeOffWords + off * 3u;
        dstWedgePts[base + 0u] = v00.x;
        dstWedgePts[base + 1u] = v00.y;
        dstWedgePts[base + 2u] = v00.z;
        dstWedgePts[base + 3u] = v10.x;
        dstWedgePts[base + 4u] = v10.y;
        dstWedgePts[base + 5u] = v10.z;
        dstWedgePts[base + 6u] = v20.x;
        dstWedgePts[base + 7u] = v20.y;
        dstWedgePts[base + 8u] = v20.z;
      } else {
        let off = dstFullOffset + 1u + 12u * threadCorner;
        let base = U.wedgeOffWords + off * 3u;
        // 12 packed Vec3s: v00 v01 v02 v03  v10 v11 v12 v13  v20 v21 v22 v23
        let vs = array<vec3<f32>, 12>(v00, v01, v02, v03, v10, v11, v12, v13, v20, v21, v22, v23);
        for (var k: u32 = 0u; k < 12u; k = k + 1u) {
          dstWedgePts[base + k * 3u + 0u] = vs[k].x;
          dstWedgePts[base + k * 3u + 1u] = vs[k].y;
          dstWedgePts[base + k * 3u + 2u] = vs[k].z;
        }
      }
    }
    workgroupBarrier();

    // ---- Ping-pong: swap buf indices (no shared-mem traffic, no
    // extra barrier). prevBuf/currBuf are uniform u32 vars; the
    // post-SAVE barrier above already guaranteed all writes to the
    // current-iter buffer are visible workgroup-wide, so iter N+1's
    // PASS 1 reads them through prevBuf (now pointing at what was
    // currBuf) without any further sync. The buffer that previously
    // held iter N-1's data becomes the new currBuf -- its stale
    // contents are about to be overwritten by PASS 1-4 anyway.
    let _tmpBuf = prevBuf;
    prevBuf = currBuf;
    currBuf = _tmpBuf;
  }
}
`;

if (typeof module !== 'undefined' && module.exports) {
  module.exports = {
    CC_GPU_LIMIT_PN_WGSL,
    CC_GPU_LIMIT_POS_WGSL,
    CC_GPU_WEDGE_ITER_WGSL,
    CC_GPU_REUSABLE_CP_WGSL,
    CcGpuLimitCtx,
    CcGpuWedgeCtx,
    CcGpuReusableCpCtx,
    gpuComputeCcLimitPN,
    gpuComputeCcLimitPositions,
    gpuComputeCcWedgeIter,
  };
}
