Skip to content

Commit e6f7120

Browse files
COMMENCINGTHESCOURGECOMMENCINGTHESCOURGE
committed
feat(compute/qef): implement watertight Marching Tetrahedra indexing pipeline
- Replace neighbor-quad fanning in mesh_assembly.wgsl with true 16-case MT triangulation case table. - Allocate and bind crossing_lut buffer in qef_pipeline.rs to track edge crossings programmatically. - Upgrade qef_solve.wgsl to resolve vertices 1-to-1 based on crossing indices, removing atomic-addition index collisions. - Add division-by-zero safeguards and clamp interpolation parameters in marching_tets.wgsl. - Integrate hydraulic host, culling, and validation rainfall adjustments in game compute pass configurations.
1 parent e0f27b0 commit e6f7120

8 files changed

Lines changed: 308 additions & 196 deletions

File tree

‎game/compute/hydraulic_host.js‎

Lines changed: 63 additions & 73 deletions
Original file line numberDiff line numberDiff line change
@@ -3,24 +3,16 @@
33
// =============================================================================
44
// Pass 1: Culling compute — reads brick_metadata, writes indirect dispatch buffer
55
// Pass 2: Solver compute — indirect dispatch from Pass 1 output
6-
//
7-
// Architecture:
8-
// Browser owns the runtime (WGSL shader + this host).
9-
// Kaggle owns the truth (material tensor calibration).
10-
// The WGSL shader and WASM fallback are both compiled artifacts from
11-
// the same hydraulic solver logic — this host loads WGSL; a Kaggle
12-
// notebook produces the WASM fallback separately.
136
// =============================================================================
147

158
const WORLD_DIM = 256;
16-
const BRICK_DIM = 8;
17-
const BRICKS_X = WORLD_DIM / BRICK_DIM; // 32
18-
const BRICKS_Y = WORLD_DIM / BRICK_DIM; // 32
19-
const LAYERS_PER_DISPATCH = 4;
20-
const BRICKS_Z_DISPATCH = WORLD_DIM / LAYERS_PER_DISPATCH; // 64
9+
const BRICK_DIM = 16;
10+
const BRICKS_X = WORLD_DIM / BRICK_DIM; // 16
11+
const BRICKS_Y = WORLD_DIM / BRICK_DIM; // 16
12+
const BRICKS_Z = WORLD_DIM / BRICK_DIM; // 16
2113

2214
const TOTAL_VOXELS = WORLD_DIM * WORLD_DIM * WORLD_DIM; // 16,777,216
23-
const TOTAL_BRICK_SLICES = BRICKS_X * BRICKS_Y * BRICKS_Z_DISPATCH; // 65,536
15+
const TOTAL_BRICK_SLICES = BRICKS_X * BRICKS_Y * BRICKS_Z; // 4,096
2416
const MAX_ACTIVE_BRICKS = TOTAL_BRICK_SLICES;
2517

2618
const BUDGET_MAX_DEFAULT = 3000;
@@ -34,27 +26,32 @@ export class HydraulicPipeline {
3426
this.device = device;
3527
this.budgetMax = budgetMax;
3628

37-
// --- SoA Voxel Buffers ---
38-
// f16 = 2 bytes/elem; vec3<f16> = 6 bytes/elem
29+
// --- SoA Voxel Buffers (u16 per voxel = 2 bytes) ---
3930
this.voxelWater = this._createStorage(TOTAL_VOXELS * 2);
31+
this.voxelWaterDst = this._createStorage(TOTAL_VOXELS * 2);
4032
this.voxelSediment = this._createStorage(TOTAL_VOXELS * 2);
41-
this.voxelPerm = this._createStorage(TOTAL_VOXELS * 6);
33+
this.voxelPermX = this._createStorage(TOTAL_VOXELS * 2);
34+
this.voxelPermY = this._createStorage(TOTAL_VOXELS * 2);
35+
this.voxelPermZ = this._createStorage(TOTAL_VOXELS * 2);
4236
this.voxelCohesion = this._createStorage(TOTAL_VOXELS * 2);
4337

44-
// --- Brick Metadata (1 u32 per 8×8×4 dispatch slice) ---
45-
this.brickMetadata = this._createStorage(TOTAL_BRICK_SLICES * 4);
38+
// --- Brick Metadata (6 channels × 16 bytes per brick) ---
39+
this.brickMetadata = this._createStorage(TOTAL_BRICK_SLICES * 6 * 16);
40+
41+
// --- Brick State for culling tracking (u32 per brick) ---
42+
this.brickState = this._createStorage(TOTAL_BRICK_SLICES * 4);
4643

4744
// --- Dispatch Chain Buffers ---
48-
// Indirect dispatch: 3 × u32 = 12 bytes, requires INDIRECT usage
45+
// Indirect dispatch: 3 × u32 = 12 bytes
4946
this.indirectBuffer = device.createBuffer({
5047
size: 12,
5148
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.INDIRECT | GPUBufferUsage.COPY_DST,
5249
});
5350

54-
// Active list: ActiveBrick = 16 bytes (u32 + 12-byte vec3<u32>)
55-
this.activeList = this._createStorage(MAX_ACTIVE_BRICKS * 16);
51+
// Active list: ActiveBrick (budgeted queue index list)
52+
this.activeList = this._createStorage(MAX_ACTIVE_BRICKS * 4);
5653

57-
// Budget counter — reset per frame by CPU write
54+
// Budget counter — reset per frame by GPU-side reset kernel
5855
this.budgetCounter = device.createBuffer({
5956
size: 4,
6057
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
@@ -83,6 +80,13 @@ export class HydraulicPipeline {
8380
this._budgetedQueue = this._createStorage(MAX_ACTIVE_BRICKS * 4); // u32
8481
this._brickPriority = this._createStorage(TOTAL_BRICK_SLICES * 4); // f32
8582

83+
// Culling schedule parameters (moistureThreshold, stabilityThreshold, emaAlpha, deadband)
84+
this.schedParamsBuffer = device.createBuffer({
85+
size: 16,
86+
usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST,
87+
});
88+
device.queue.writeBuffer(this.schedParamsBuffer, 0, new Float32Array([0.1, 0.5, 0.3, 0.05]));
89+
8690
// Camera position uniform for LOD culling (vec3<f32> = 12 bytes)
8791
this._cameraPos = new Float32Array([128.0, 64.0, 128.0]);
8892
this.cameraBuffer = device.createBuffer({
@@ -100,15 +104,15 @@ export class HydraulicPipeline {
100104
// Initialization — must be called after WGSL sources are fetched
101105
// ==========================================================================
102106

103-
async init(pass1WGSL, pass2WGSL, metaDispatchWGSL) {
107+
async init(pass1WGSL, pass2WGSL, metaDispatchWGSL, editMinBuffer, editMaxBuffer) {
104108
const device = this.device;
105109

106110
const pass1Module = device.createShaderModule({ code: pass1WGSL });
107111
const pass2Module = device.createShaderModule({ code: pass2WGSL });
108112

109113
this.pass1Pipeline = device.createComputePipeline({
110114
layout: 'auto',
111-
compute: { module: pass1Module, entryPoint: 'cull_active_bricks' },
115+
compute: { module: pass1Module, entryPoint: 'culling_pass' },
112116
});
113117

114118
// --- Reset Queue Pipeline ---
@@ -136,11 +140,12 @@ export class HydraulicPipeline {
136140
this._metaBindGroup = device.createBindGroup({
137141
layout: this._metaPipeline.getBindGroupLayout(0),
138142
entries: [
139-
{ binding: 0, resource: { buffer: this.indirectBuffer }}, // compacted_queue
143+
{ binding: 0, resource: { buffer: this.activeList }}, // compacted_queue
140144
{ binding: 1, resource: { buffer: this.budgetCounter }}, // queue_count
141145
{ binding: 2, resource: { buffer: this._brickPriority }}, // brick_priority
142146
{ binding: 3, resource: { buffer: this._budgetedQueue }}, // budgeted_queue
143147
{ binding: 4, resource: { buffer: this.indirectBuffer }}, // dispatch_args (reuse mem)
148+
{ binding: 5, resource: { buffer: this._metaUniformBuffer }}, // meta_params
144149
],
145150
});
146151
// Meta params uniform
@@ -153,60 +158,57 @@ export class HydraulicPipeline {
153158

154159
this.pass2Pipeline = device.createComputePipeline({
155160
layout: 'auto',
156-
compute: { module: pass2Module, entryPoint: 'hydraulic_solver' },
161+
compute: { module: pass2Module, entryPoint: 'advection_pass' },
157162
});
158163

159164
// --- Pass 1 Bind Group ---
160-
// bindings: 0=meta_buffer, 1=brick_state, 2=raw_queue, 3=queue_count,
161-
// 4=dispatch_args, 5=sched_params, 6=camera_pos
162165
this.pass1BindGroup = device.createBindGroup({
163166
layout: this.pass1Pipeline.getBindGroupLayout(0),
164167
entries: [
165168
{ binding: 0, resource: { buffer: this.brickMetadata }},
166-
{ binding: 1, resource: { buffer: this.indirectBuffer }},
169+
{ binding: 1, resource: { buffer: this.brickState }},
167170
{ binding: 2, resource: { buffer: this.activeList }},
168171
{ binding: 3, resource: { buffer: this.budgetCounter }},
169-
{ binding: 6, resource: { buffer: this.cameraBuffer }},
172+
{ binding: 4, resource: { buffer: this.schedParamsBuffer }},
173+
{ binding: 5, resource: { buffer: this.cameraBuffer }},
170174
],
171175
});
172176

173177
// --- Pass 2 Bind Groups ---
174178
this.pass2BindGroup0 = device.createBindGroup({
175179
layout: this.pass2Pipeline.getBindGroupLayout(0),
176180
entries: [
177-
{ binding: 0, resource: { buffer: this.voxelWater }},
178-
{ binding: 1, resource: { buffer: this.voxelSediment }},
179-
{ binding: 2, resource: { buffer: this.voxelPerm }},
180-
{ binding: 3, resource: { buffer: this.voxelCohesion }},
181-
{ binding: 4, resource: { buffer: this.brickMetadata }},
181+
{ binding: 0, resource: { buffer: this.brickMetadata }},
182+
{ binding: 1, resource: { buffer: this.voxelWater }},
183+
{ binding: 2, resource: { buffer: this.voxelWaterDst }},
184+
{ binding: 3, resource: { buffer: this.voxelPermX }},
185+
{ binding: 4, resource: { buffer: this.voxelPermY }},
186+
{ binding: 5, resource: { buffer: this.voxelPermZ }},
187+
{ binding: 10, resource: { buffer: editMinBuffer }},
188+
{ binding: 11, resource: { buffer: editMaxBuffer }},
182189
],
183190
});
184191

185192
this.pass2BindGroup1 = device.createBindGroup({
186193
layout: this.pass2Pipeline.getBindGroupLayout(1),
187194
entries: [
188-
{ binding: 0, resource: { buffer: this.activeList }},
189-
{ binding: 1, resource: { buffer: this.indirectBuffer }},
190-
],
191-
});
192-
193-
this.pass2BindGroup2 = device.createBindGroup({
194-
layout: this.pass2Pipeline.getBindGroupLayout(2),
195-
entries: [
196-
{ binding: 0, resource: { buffer: this.uniformBuffer }},
195+
{ binding: 0, resource: { buffer: this._budgetedQueue }}, // compacted_queue
197196
],
198197
});
199198
}
200199

201200
// ==========================================================================
202-
// Upload initial world state from material tensor (Kaggle output or procedural)
201+
// Upload initial world state from material tensor
203202
// ==========================================================================
204203

205-
uploadInitialState({ water, sediment, permeability, cohesion, metadata }) {
204+
uploadInitialState({ water, sediment, permX, permY, permZ, cohesion, metadata }) {
206205
const q = this.device.queue;
207206
q.writeBuffer(this.voxelWater, 0, water);
207+
q.writeBuffer(this.voxelWaterDst, 0, water);
208208
q.writeBuffer(this.voxelSediment, 0, sediment);
209-
q.writeBuffer(this.voxelPerm, 0, permeability);
209+
q.writeBuffer(this.voxelPermX, 0, permX);
210+
q.writeBuffer(this.voxelPermY, 0, permY);
211+
q.writeBuffer(this.voxelPermZ, 0, permZ);
210212
q.writeBuffer(this.voxelCohesion, 0, cohesion);
211213
q.writeBuffer(this.brickMetadata, 0, metadata);
212214
}
@@ -228,7 +230,7 @@ export class HydraulicPipeline {
228230
queue.writeBuffer(this.cameraBuffer, 0, this._cameraPos);
229231

230232
// ========================================================================
231-
// Pass 0: GPU-side queue reset — replaces host writeBuffer
233+
// Pass 0: GPU-side queue reset
232234
// ========================================================================
233235
{
234236
const pass = cmd.beginComputePass();
@@ -239,13 +241,12 @@ export class HydraulicPipeline {
239241
}
240242

241243
// ========================================================================
242-
// Pass 1: Culling — determine which bricks are active
244+
// Pass 1: Culling — determine active bricks
243245
// ========================================================================
244246
{
245247
const pass = cmd.beginComputePass();
246248
pass.setPipeline(this.pass1Pipeline);
247249
pass.setBindGroup(0, this.pass1BindGroup);
248-
pass.setBindGroup(1, this._cullingUniformBindGroup);
249250

250251
const workgroups = Math.ceil(TOTAL_BRICK_SLICES / 64);
251252
pass.dispatchWorkgroups(workgroups, 1, 1);
@@ -262,26 +263,28 @@ export class HydraulicPipeline {
262263
const pass = cmd.beginComputePass();
263264
pass.setPipeline(this._metaPipeline);
264265
pass.setBindGroup(0, this._metaBindGroup);
265-
// Dispatch enough WGs to cover MAX_ACTIVE_BRICKS (65536) at 256 threads ea
266-
pass.dispatchWorkgroups(256, 1, 1);
266+
pass.dispatchWorkgroups(16, 1, 1); // 16 workgroups * 256 threads covers 4096 bricks
267267
pass.end();
268268
}
269269

270270
// ========================================================================
271-
// Pass 2: Hydraulic solver — only dispatched bricks run
271+
// Pass 2: Hydraulic solver — only active bricks run
272272
// ========================================================================
273273
{
274274
const pass = cmd.beginComputePass();
275275
pass.setPipeline(this.pass2Pipeline);
276276
pass.setBindGroup(0, this.pass2BindGroup0);
277277
pass.setBindGroup(1, this.pass2BindGroup1);
278-
pass.setBindGroup(2, this.pass2BindGroup2);
279278

280-
// GPU driver reads indirectBuffer.x → dispatches exactly that many
281279
pass.dispatchWorkgroupsIndirect(this.indirectBuffer, 0);
282280
pass.end();
283281
}
284282

283+
// ========================================================================
284+
// Pass 3: Water Ping-Pong Buffer Swap
285+
// ========================================================================
286+
cmd.copyBufferToBuffer(this.voxelWaterDst, 0, this.voxelWater, 0, TOTAL_VOXELS * 2);
287+
285288
queue.submit([cmd.finish()]);
286289
}
287290

@@ -304,7 +307,7 @@ export class HydraulicPipeline {
304307
}
305308

306309
// ==========================================================================
307-
// Update budget uniform (no kernel recompile needed)
310+
// Update budget uniform
308311
// ==========================================================================
309312

310313
setBudgetMax(newBudget) {
@@ -320,36 +323,23 @@ export class HydraulicPipeline {
320323
_createStorage(size) {
321324
return this.device.createBuffer({
322325
size,
323-
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
326+
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST | GPUBufferUsage.COPY_SRC,
324327
});
325328
}
326-
327-
// Culling uniform bind group (reused)
328-
get _cullingUniformBindGroup() {
329-
if (!this.__cullingBG) {
330-
this.__cullingBG = this.device.createBindGroup({
331-
layout: this.pass1Pipeline.getBindGroupLayout(1),
332-
entries: [
333-
{ binding: 0, resource: { buffer: this.uniformBuffer }},
334-
],
335-
});
336-
}
337-
return this.__cullingBG;
338-
}
339329
}
340330

341331
// =============================================================================
342-
// Bootstrap helper — loads WGSL sources and initializes the pipeline
332+
// Bootstrap helper
343333
// =============================================================================
344334

345-
export async function createHydraulicPipeline(device, budgetMax) {
335+
export async function createHydraulicPipeline(device, budgetMax, editMinBuffer, editMaxBuffer) {
346336
const [pass1WGSL, pass2WGSL, metaDispatchWGSL] = await Promise.all([
347337
fetch('compute/pass1_culling.wgsl').then(r => r.text()),
348338
fetch('compute/pass2_solver.wgsl').then(r => r.text()),
349339
fetch('compute/meta_dispatch.wgsl').then(r => r.text()),
350340
]);
351341

352342
const pipeline = new HydraulicPipeline(device, budgetMax);
353-
await pipeline.init(pass1WGSL, pass2WGSL, metaDispatchWGSL);
343+
await pipeline.init(pass1WGSL, pass2WGSL, metaDispatchWGSL, editMinBuffer, editMaxBuffer);
354344
return pipeline;
355345
}

‎game/compute/pass1_culling.wgsl‎

Lines changed: 2 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -16,9 +16,8 @@ struct DispatchArgs { count_x: u32, count_y: u32, count_z: u32 }
1616
@group(0) @binding(1) var<storage, read_write> brick_state: array<u32>;
1717
@group(0) @binding(2) var<storage, read_write> raw_queue: array<u32>;
1818
@group(0) @binding(3) var<storage, read_write> queue_count: atomic<u32>;
19-
@group(0) @binding(4) var<uniform> dispatch_args: DispatchArgs;
20-
@group(0) @binding(5) var<uniform> sched_params: vec4<f32>;
21-
@group(0) @binding(6) var<uniform> camera_pos: vec3<f32>;
19+
@group(0) @binding(4) var<uniform> sched_params: vec4<f32>;
20+
@group(0) @binding(5) var<uniform> camera_pos: vec3<f32>;
2221

2322
@compute @workgroup_size(64)
2423
fn culling_pass(@builtin(global_invocation_id) gid: vec3<u32>) {

‎game/compute/pass2_solver.wgsl‎

Lines changed: 1 addition & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -79,11 +79,7 @@ const GRAVITY: f32 = 9.81;
7979
@group(0) @binding(4) var<storage, read> perm_y_u16: array<F16>;
8080
@group(0) @binding(5) var<storage, read> perm_z_u16: array<F16>;
8181

82-
@group(1) @binding(0) var<storage, read> active_list: array<ActiveBrick>;
83-
@group(1) @binding(1) var<storage, read> dispatch_indirect: DispatchIndirect;
84-
85-
// ── Compacted queue from scheduling pass ──
86-
@group(1) @binding(2) var<storage, read> compacted_queue: array<u32>;
82+
@group(1) @binding(0) var<storage, read> compacted_queue: array<u32>;
8783

8884
// =============================================================================
8985
// ── Phase 6A: Dynamic Range Shadow Buffer Decode ──

0 commit comments

Comments
 (0)