fix: phase5 & phase6 GPU compute audit — sync barriers, buffer structs, shader guards

game/compute/phase5_extractor/phase5_draw.js

@@ -23,6 +23,14 @@ export class Phase5Draw {
       usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST | GPUBufferUsage.COPY_SRC,
     });
 
+    // ISO threshold uniform buffer (1 f32 = 4 bytes)
+    this.isoThresholdBuffer = device.createBuffer({
+      size: 4,
+      usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST,
+    });
+    // Default ISO threshold = 0.5
+    device.queue.writeBuffer(this.isoThresholdBuffer, 0, new Float32Array([0.5]));
+
     // Indirect draw args: [indexCount, instanceCount, firstIndex, baseVertex, firstInstance]
     this.drawArgsBuffer = device.createBuffer({
       size: 20,
@@ -59,15 +67,19 @@ export class Phase5Draw {
     });
   }
 
-  _createTopologyBG(qefBuffer, hermiteBuffer, isoThreshold) {
+  _createTopologyBG(qefBuffer, hermiteBuffer, isoThresholdValue) {
+    // Update ISO threshold buffer if a value is provided
+    if (isoThresholdValue !== undefined) {
+      this.device.queue.writeBuffer(this.isoThresholdBuffer, 0, new Float32Array([isoThresholdValue]));
+    }
     return this.device.createBindGroup({
       layout: this.pipelines.topology.getBindGroupLayout(0),
       entries: [
         { binding: 0, resource: { buffer: qefBuffer } },
         { binding: 1, resource: { buffer: hermiteBuffer } },
         { binding: 2, resource: { buffer: this.indexBuffer } },
         { binding: 3, resource: { buffer: this.indexCountBuffer } },
-        { binding: 4, resource: { buffer: isoThreshold } },
+        { binding: 4, resource: { buffer: this.isoThresholdBuffer } },
       ],
     });
   }
@@ -84,24 +96,28 @@ export class Phase5Draw {
 
   /**
    * Full mesh build: topology generation + LOD stitching + draw args.
+   * 
+   * Topology: generates faces for all 255³ cells.
+   * LOD stitching: corrects seams between LOD levels.
    */
-  async buildMesh(qefBuffer, hermiteBuffer, lodBuffer, isoThresholdValue) {
+  async buildMesh(qefBuffer, hermiteBuffer, lodBuffer, isoThresholdValue = 0.5) {
     const device = this.device;
     const encoder = device.createCommandEncoder();
 
     // Reset index count
     device.queue.writeBuffer(this.indexCountBuffer, 0, new Uint32Array([0]));
 
-    // --- Pass 1: Topology generation ---
+    // --- Pass 1: Topology generation (full 3D: 255³ cells) ---
     {
       const pass = encoder.beginComputePass();
       pass.setPipeline(this.pipelines.topology);
       pass.setBindGroup(0, this._createTopologyBG(qefBuffer, hermiteBuffer, isoThresholdValue));
-      pass.dispatchWorkgroups(32, 32, 1);  // 255 × 255 cells in XY
+      // Dispatch: 32×32×32 workgroups = 1024³ threads covering 255³ cells
+      pass.dispatchWorkgroups(32, 32, 32);
       pass.end();
     }
 
-    // --- Pass 2: LOD stitching ---
+    // --- Pass 2: LOD stitching (full 3D: 255³ cells) ---
     {
       const pass = encoder.beginComputePass();
       pass.setPipeline(this.pipelines.stitch);
@@ -113,8 +129,7 @@ export class Phase5Draw {
     // --- Update indirect draw args ---
     // copy indexCount → drawArgs[0] (indexCount)
     encoder.copyBufferToBuffer(this.indexCountBuffer, 0, this.drawArgsBuffer, 0, 4);
-    // Set instanceCount = 1 (at offset 4)
-    // In a separate writeBuffer:
+    // Set instanceCount = 1, firstIndex = 0, baseVertex = 0, firstInstance = 0
     device.queue.writeBuffer(this.drawArgsBuffer, 4, new Uint32Array([1, 0, 0, 0]));
 
     device.queue.submit([encoder.finish()]);
@@ -149,7 +164,7 @@ export class Phase5Draw {
   }
 
   destroy() {
-    const bufs = ['indexBuffer', 'indexCountBuffer', 'drawArgsBuffer', 'readbackBuffer'];
+    const bufs = ['indexBuffer', 'indexCountBuffer', 'isoThresholdBuffer', 'drawArgsBuffer', 'readbackBuffer'];
     for (const key of bufs) {
       if (this[key]) this[key].destroy();
     }

game/compute/phase5_extractor/phase5_host.js

@@ -6,12 +6,22 @@
  * Writes GPU-visible dual vertex buffer for rendering.
  */
 
+const CHANNELS = {
+  DENSITY: 0,
+  SEDIMENT: 1,
+  PERM_X: 2,
+  PERM_Y: 3,
+  PERM_Z: 4,
+  COHESION: 5,
+};
+
 export class Phase5Extractor {
   constructor(device, gridSize = 256) {
     this.device = device;
     this.gridSize = gridSize;
     this.cellCount = (gridSize - 1) ** 3;    // 255³
     this.vertexCount = (gridSize + 1) ** 3;   // 257³
+    this.CHANNELS = CHANNELS;
 
     this._createBuffers();
     this.pipelines = {};
@@ -91,6 +101,13 @@ export class Phase5Extractor {
   }
 
   _createHermiteBG(metaBuffer, densityBuf, cohesionBuf, permXBuf) {
+    // Validate buffers exist and are GPU buffers
+    if (!densityBuf || !cohesionBuf || !permXBuf || !metaBuffer) {
+      throw new Error('Phase5Extractor: hermite bind group missing required buffers (density, cohesion, permX, meta)');
+    }
+    if (!densityBuf.size || !cohesionBuf.size || !permXBuf.size || !metaBuffer.size) {
+      throw new Error('Phase5Extractor: hermite buffers have zero size');
+    }
     return this.device.createBindGroup({
       layout: this.pipelines.hermite.getBindGroupLayout(0),
       entries: [
@@ -139,12 +156,13 @@ export class Phase5Extractor {
     });
   }
 
-  async fullExtract(metaBuffer, channelBuffers, brickMetaBuffer) {
+  async fullExtract(metaBuffer, channelBuffers, brickMetaBuffer, qefParams = { tolerance: 0.3, weightThreshold: 0.01 }) {
     const device = this.device;
     const encoder = device.createCommandEncoder();
 
-    // Set QEF params
-    device.queue.writeBuffer(this.qefParamsBuffer, 0, new Float32Array([0.3, 0.01]));
+    // Set QEF params (configurable per-call)
+    const paramData = new Float32Array([qefParams.tolerance, qefParams.weightThreshold]);
+    device.queue.writeBuffer(this.qefParamsBuffer, 0, paramData);
 
     // Pass 1: Hermite data generation (32×32×32 workgroups = 256³ vertices)
     {
@@ -181,6 +199,7 @@ export class Phase5Extractor {
     const copyEncoder = device.createCommandEncoder();
     copyEncoder.copyBufferToBuffer(this.vertexBuffer, 0, this.prevVertexBuffer, 0, this.cellCount * 12);
     device.queue.submit([copyEncoder.finish()]);
+    await device.queue.onSubmittedWorkDone();
   }
 
   async incrementalExtract(brickMetaBuffer) {
@@ -207,6 +226,7 @@ export class Phase5Extractor {
     const readEncoder = device.createCommandEncoder();
     readEncoder.copyBufferToBuffer(this.deltaCountBuffer, 0, readback, 0, 4);
     device.queue.submit([readEncoder.finish()]);
+    await device.queue.onSubmittedWorkDone();
 
     await readback.mapAsync(GPUMapMode.READ);
     const count = new Uint32Array(readback.getMappedRange())[0];
@@ -217,6 +237,7 @@ export class Phase5Extractor {
     const copyEncoder = device.createCommandEncoder();
     copyEncoder.copyBufferToBuffer(this.vertexBuffer, 0, this.prevVertexBuffer, 0, this.cellCount * 12);
     device.queue.submit([copyEncoder.finish()]);
+    await device.queue.onSubmittedWorkDone();
 
     return count;
   }

game/compute/phase5_extractor/tiled_extractor.js

@@ -120,6 +120,12 @@ export class TiledExtractor {
       size: 8,
       usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST,
     });
+
+    // Tile offset uniform (vec3<u32> = 12 bytes, padded to 16)
+    this.tileOffsetBuffer = this.device.createBuffer({
+      size: 16,
+      usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST,
+    });
   }
 
   async init(wgslSources) {
@@ -163,6 +169,10 @@ export class TiledExtractor {
     for (let tz = 0; tz < this.tilesPerDim; tz++) {
       for (let ty = 0; ty < this.tilesPerDim; ty++) {
         for (let tx = 0; tx < this.tilesPerDim; tx++) {
+          // Tile loop offset for global buffer writes
+          const tileOffsetData = new Uint32Array([vx0, vy0, vz0, 0]); // padding
+          queue.writeBuffer(this.tileOffsetBuffer, 0, tileOffsetData);
+
           const encoder = device.createCommandEncoder();
 
           // Compute tile bounds in vertex space
@@ -242,13 +252,14 @@ export class TiledExtractor {
   }
 
   _createQEFBG_Tiled(tileX, tileY, tileZ) {
-    // TODO: Add tile_offset uniform binding once WGSL accepts it
+    // Now includes tile_offset uniform binding
     return this.device.createBindGroup({
       layout: this.pipelines.qef.getBindGroupLayout(0),
       entries: [
         { binding: 0, resource: { buffer: this.tileHermiteBuffer } },
         { binding: 1, resource: { buffer: this.vertexBuffer } },
         { binding: 2, resource: { buffer: this.qefParamsBuffer } },
+        { binding: 3, resource: { buffer: this.tileOffsetBuffer } },
       ],
     });
   }
@@ -261,16 +272,25 @@ export class TiledExtractor {
         { binding: 1, resource: { buffer: permXBuf } },
         { binding: 2, resource: { buffer: this.lodBuffer } },
         { binding: 3, resource: { buffer: metaBuffer } },
+        { binding: 4, resource: { buffer: this.tileOffsetBuffer } },
       ],
     });
   }
 
   getVertexBuffer() { return this.vertexBuffer; }
   getLODBuffer() { return this.lodBuffer; }
+
+  destroy() {
+    const bufs = ['vertexBuffer', 'prevVertexBuffer', 'lodBuffer', 'tileHermiteBuffer',
+                  'deltaBuffer', 'deltaCountBuffer', 'qefParamsBuffer', 'tileOffsetBuffer'];
+    for (const key of bufs) {
+      if (this[key]) this[key].destroy();
+    }
+  }
 }
 
 /* 
- * WGSL SHADER MODIFICATION REQUIRED
+ * WGSL SHADER MODIFICATIONS REQUIRED
  * 
  * The hermite and qef shaders currently compute global vertex indices
  * from global_invocation_id. For tiled extraction, they need an
@@ -283,11 +303,15 @@ export class TiledExtractor {
  *   let vx = gid.x + 1u + tile_offset.x;
  *   let vy = gid.y + 1u + tile_offset.y;
  *   let vz = gid.z + 1u + tile_offset.z;
+ *   let vertex_idx = vx + vy * GRID_SIZE + vz * GRID_SIZE * GRID_SIZE;
  *   
- * Without this uniform, the shaders must be modified or the tiled
- * approach requires one bind group per tile position (impractical).
+ *   // In qef_solve, similarly adjust cell indices:
+ *   let cx = gid.x + tile_offset.x;
+ *   let cy = gid.y + tile_offset.y;
+ *   let cz = gid.z + tile_offset.z;
  * 
- * For a minimal first pass: use the original full-hermite allocation
- * and only tile the QEF+LOD passes (which are smaller). That drops
- * peak from 942MB to ~667MB with less shader modification.
+ * Binding locations:
+ *   hermite: binding(5) = tile_offset
+ *   qef: binding(3) = tile_offset
+ *   lod: binding(4) = tile_offset
  */

game/compute/phase6_edit/phase6_edit.wgsl

@@ -46,6 +46,13 @@ const VOXELS_PER_BRICK: u32 = 4096u;
 // ── Helper: Encode with Range Expansion Detection ─────────────
 fn encode_with_expansion(ch: u32, brick_idx: u32, local_idx: u32, val: f32, dst: ptr<function, array<F16>>) -> bool {
   let meta = brick_meta[brick_idx * 6u + ch];
+
+  // Guard against division by zero
+  if (meta.y < 1e-6) {
+    (*dst)[brick_idx * VOXELS_PER_BRICK + local_idx] = f16_encode(clamp(val, 0.0, 1.0));
+    return false; // No expansion needed if range is degenerate
+  }
+
   let norm = (val - meta.x) / meta.y;
 
   let needs_expand_min = (val < meta.x);
@@ -108,10 +115,12 @@ fn inject_pass(@builtin(global_invocation_id) gid: vec3<u32>) {
 
         if (cmd.material_type == 0u) { // Carve → Air
           let t = smoothstep(cmd.radius, cmd.radius * (1.0 - cmd.falloff), d);
-          new_density = mix(0.0, 1.0, t);
+          // t=1 at outer edge (should stay solid), t=0 at center (should be air)
+          // So density goes from 1.0 (outer) to 0.0 (center) as we carve inward
+          new_density = mix(0.0, 1.0, t);  // Correct: 1.0 * t + 0.0 * (1-t)
           new_cohesion = mix(0.0, 1.0, t);
           new_perm = 1.0;
-          clear_water = (t > 0.5); // Clear water in fully carved voxels
+          clear_water = (t < 0.5); // Clear water in carved voxels (center)
         } else if (cmd.material_type == 1u) { // Inject Ore
           let t = smoothstep(cmd.radius * 0.5, cmd.radius, d);
           new_density = mix(0.95, 0.5, t);
@@ -124,17 +133,19 @@ fn inject_pass(@builtin(global_invocation_id) gid: vec3<u32>) {
         }
 
         // Encode with range expansion detection
+        // Note: atomicMin/atomicMax expect u32 pointers, so ensure the buffers are u32
         if (encode_with_expansion(0u, brick_idx, local_idx, new_density, &density_u16)) {
-          atomicMin(&edit_min_buffer[brick_idx * 6u + 0u], F16(new_density));
-          atomicMax(&edit_max_buffer[brick_idx * 6u + 0u], F16(new_density));
+          // Store expanded bounds as u32 bit patterns representing F16 values
+          atomicMin(&edit_min_buffer[brick_idx * 6u + 0u], f16_bits(new_density));
+          atomicMax(&edit_max_buffer[brick_idx * 6u + 0u], f16_bits(new_density));
         }
         if (encode_with_expansion(1u, brick_idx, local_idx, new_cohesion, &cohesion_u16)) {
-          atomicMin(&edit_min_buffer[brick_idx * 6u + 1u], F16(new_cohesion));
-          atomicMax(&edit_max_buffer[brick_idx * 6u + 1u], F16(new_cohesion));
+          atomicMin(&edit_min_buffer[brick_idx * 6u + 1u], f16_bits(new_cohesion));
+          atomicMax(&edit_max_buffer[brick_idx * 6u + 1u], f16_bits(new_cohesion));
         }
         if (encode_with_expansion(2u, brick_idx, local_idx, new_perm, &perm_x_u16)) {
-          atomicMin(&edit_min_buffer[brick_idx * 6u + 2u], F16(new_perm));
-          atomicMax(&edit_max_buffer[brick_idx * 6u + 2u], F16(new_perm));
+          atomicMin(&edit_min_buffer[brick_idx * 6u + 2u], f16_bits(new_perm));
+          atomicMax(&edit_max_buffer[brick_idx * 6u + 2u], f16_bits(new_perm));
         }
 
         // Clear water in carved voids (prevent physics explosion)

game/compute/phase6_edit/phase6_host.js

@@ -13,8 +13,8 @@ export class EditManager {
     this.nextSlot = 0;
     this.pendingCount = 0;
 
-    // Command descriptor
-    this.commandByteSize = 32; // sizeof(EditCommand) padded
+    // Command descriptor: center(12B) + radius(4B) + materialType(4B) + falloff(4B) = 24 bytes
+    this.commandByteSize = 24;
 
     // Ring buffer for edit commands (GPU-visible)
     this.editBuffer = device.createBuffer({
@@ -163,7 +163,17 @@ export class EditManager {
     pass.dispatchWorkgroups(wgX, wgY, 1);
     pass.end();
 
-    // Reset counter for next frame
+    // NOTE: Do NOT reset editCountBuffer here. It must be reset AFTER
+    // this encoder finishes on the GPU. Use resetEditCount() as a separate
+    // GPU pass or call it after await device.queue.onSubmittedWorkDone().
+  }
+
+  /**
+   * Reset edit counter. Must be called AFTER applyEdits() GPU work is done.
+   * Can be called on a separate command encoder or via writeBuffer after
+   * GPU completion.
+   */
+  resetEditCount() {
     this.device.queue.writeBuffer(this.editCountBuffer, 0, new Uint32Array([0]));
     this.pendingCount = 0;
   }

qef_extraction/density_field.wgsl

@@ -0,0 +1,44 @@
+// density_field.wgsl
+// Kernel 1: Convert particle positions to 64³ density field.
+// Reads spatial_hash.wgsl grid, writes density scalar per cell.
+//
+// Dispatch: (64, 64, 1) workgroups of (1, 1, 64) — one thread per z-slice.
+
+struct DensityParams {
+    grid_dim: u32,            // 64
+    max_particles_per_cell: u32,
+    particle_radius: f32,     // for density falloff
+};
+
+@group(0) @binding(0) var<storage, read> grid_heads: array<atomic<i32>>;
+@group(0) @binding(1) var<storage, read> grid_next: array<atomic<i32>>;
+@group(0) @binding(2) var<storage, read> particle_positions: array<vec3<f32>>;
+@group(0) @binding(3) var<storage, read_write> density_field: array<f32>; // 64³ = 262,144
+@group(0) @binding(4) var<uniform> params: DensityParams;
+
+@compute @workgroup_size(8, 8, 1)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+    let cx = gid.x;
+    let cy = gid.y;
+    let cz = gid.z;
+
+    if (cx >= params.grid_dim || cy >= params.grid_dim || cz >= params.grid_dim) {
+        return;
+    }
+
+    let cell_idx = cx + cy * params.grid_dim + cz * params.grid_dim * params.grid_dim;
+
+    // Count particles in this cell by walking linked list
+    var count = 0u;
+    var curr = atomicLoad(&grid_heads[cell_idx]);
+
+    while (curr >= 0 && count < params.max_particles_per_cell) {
+        count++;
+        let n_idx = u32(curr);
+        curr = atomicLoad(&grid_next[n_idx]);
+    }
+
+    // Normalize: density = count / max (capped at 1.0)
+    let density = f32(count) / f32(params.max_particles_per_cell);
+    density_field[cell_idx] = density;
+}