struct Uniforms { numTriangles : u32; gridSize : u32; pad1 : u32; pad2 : u32; bbMin : vec3; // offset(16) bbMax : vec3; // offset(32) }; struct Dbg { offsetCounter : atomic; pad0 : u32; pad1 : u32; pad2 : u32; value0 : u32; value1 : u32; value2 : u32; value3 : u32; value_f32_0 : f32; value_f32_1 : f32; value_f32_2 : f32; value_f32_3 : f32; }; struct F32s { values : [[stride(4)]] array; }; struct U32s { values : [[stride(4)]] array; }; struct I32s { values : [[stride(4)]] array; }; struct AU32s { values : [[stride(4)]] array>; }; struct AI32s { values : [[stride(4)]] array>; }; // IN [[binding(0), group(0)]] var uniforms : Uniforms; [[binding(10), group(0)]] var indices : U32s; [[binding(11), group(0)]] var positions : F32s; // OUT [[binding(20), group(0)]] var counters : AU32s; [[binding(21), group(0)]] var LUT : AI32s; // DEBUG [[binding(50), group(0)]] var dbg : Dbg; fn toVoxelPos(position : vec3) -> vec3{ var bbMin = vec3(uniforms.bbMin.x, uniforms.bbMin.y, uniforms.bbMin.z); var bbMax = vec3(uniforms.bbMax.x, uniforms.bbMax.y, uniforms.bbMax.z); var bbSize = bbMax - bbMin; var cubeSize = max(max(bbSize.x, bbSize.y), bbSize.z); var gridSize = f32(uniforms.gridSize); var gx = gridSize * (position.x - uniforms.bbMin.x) / cubeSize; var gy = gridSize * (position.y - uniforms.bbMin.y) / cubeSize; var gz = gridSize * (position.z - uniforms.bbMin.z) / cubeSize; return vec3(gx, gy, gz); } fn toIndex1D(gridSize : u32, voxelPos : vec3) -> u32{ var icoord = vec3(voxelPos); return icoord.x + gridSize * icoord.y + gridSize * gridSize * icoord.z; } fn toIndex3D(gridSize : u32, index : u32) -> vec3{ var z = index / (gridSize * gridSize); var y = (index - gridSize * gridSize * z) / gridSize; var x = index % gridSize; return vec3(x, y, z); } fn loadPosition(vertexIndex : u32) -> vec3 { var position = vec3( positions.values[3u * vertexIndex + 0u], positions.values[3u * vertexIndex + 1u], positions.values[3u * vertexIndex + 2u], ); return position; }; fn doIgnore(){ var g42 = uniforms.numTriangles; var kj6 = dbg.value1; var b53 = atomicLoad(&counters.values[0]); var rwg = indices.values[0]; var rb5 = positions.values[0]; var g55 = atomicLoad(&LUT.values[0]); } [[stage(compute), workgroup_size(128)]] fn main_count([[builtin(global_invocation_id)]] GlobalInvocationID : vec3) { var triangleIndex = GlobalInvocationID.x; if(triangleIndex >= uniforms.numTriangles){ return; } doIgnore(); var i0 = indices.values[3u * triangleIndex + 0u]; var i1 = indices.values[3u * triangleIndex + 1u]; var i2 = indices.values[3u * triangleIndex + 2u]; var p0 = loadPosition(i0); var p1 = loadPosition(i1); var p2 = loadPosition(i2); var center = (p0 + p1 + p2) / 3.0; var voxelPos = toVoxelPos(center); var voxelIndex = toIndex1D(uniforms.gridSize, voxelPos); var acefg = atomicAdd(&counters.values[voxelIndex], 1u); if(triangleIndex == 0u){ dbg.value0 = uniforms.gridSize; // dbg.value1 = voxelPos.y; // dbg.value2 = voxelPos.z; dbg.value_f32_0 = center.x; dbg.value_f32_1 = center.y; dbg.value_f32_2 = center.z; } } [[stage(compute), workgroup_size(128)]] fn main_create_lut([[builtin(global_invocation_id)]] GlobalInvocationID : vec3) { var voxelIndex = GlobalInvocationID.x; doIgnore(); var maxVoxels = uniforms.gridSize * uniforms.gridSize * uniforms.gridSize; if(voxelIndex >= maxVoxels){ return; } var numTriangles = atomicLoad(&counters.values[voxelIndex]); var offset = -1; if(numTriangles > 0u){ offset = i32(atomicAdd(&dbg.offsetCounter, numTriangles)); } atomicStore(&LUT.values[voxelIndex], offset); } [[stage(compute), workgroup_size(128)]] fn main_sort_triangles([[builtin(global_invocation_id)]] GlobalInvocationID : vec3) { var triangleIndex = GlobalInvocationID.x; doIgnore(); if(triangleIndex >= uniforms.numTriangles){ return; } var i0 = indices.values[3u * triangleIndex + 0u]; var i1 = indices.values[3u * triangleIndex + 1u]; var i2 = indices.values[3u * triangleIndex + 2u]; var p0 = loadPosition(i0); var p1 = loadPosition(i1); var p2 = loadPosition(i2); var center = (p0 + p1 + p2) / 3.0; var voxelPos = toVoxelPos(center); var voxelIndex = toIndex1D(uniforms.gridSize, voxelPos); // var triangleOffset = LUT.values[voxelIndex]; var triangleOffset = atomicAdd(&LUT.values[voxelIndex], 1); }