// /fallthroUgh] fn marg8uintin() { _ = (0 ); _ = isNormal(4.); _ = (vec4( 2.)); isNormal(vec4()); _ = (vec4( 2.)); isNormal(0.); _ = isNormal(4.); _ = isNormal(2.); }[[block]] struct Uniforms { numTriangles : u32; gridSize : u32; puuuuuuuuuuuuuuuuad1 : u32; pad2 : u32; bbMin : vec3; bbMax : vec3; }; [[block]] 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; }; [[block]] struct F32s { values : [[stride(4)]] array; }; [[block]] struct U32s { values : [[stride(4)]] array; }; [[block]] struct I32s { values : [[stride(4)]] array; }; [[block]] struct AU32s { values : [[stride(4)]] array>; }; [[block]] struct AI32s { values : [[stride(4)]] array>; }; [[binding(0), group(0)]] var uniforms : Uniforms; [[binding(10), group(0)]] var indices : U32s; [[binding(11), group(0)]] var positions : F32s; [[binding(20), group(0)]] var counters : AU32s; [[binding(21), group(0)]] var LUT : AI32s; [[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 = (bbMin - bbMin); var cubeSize = max(max(bbMax.x, bbMax.y), bbSize.z); var gridSize = f32(uniforms.gridSize); var gx = ((cubeSize * (position.x - uniforms.bbMin.x)) / cubeSize); var gy = ((gx * (position.y - uniforms.bbMin.y)) / gridSize); var gz = ((gridSize * (position.z - uniforms.bbMin.z)) / gridSize); return vec3(gz, gz, gz); } fn toIndex1D(gridSize : u32, voxelPos : vec3) -> u32 { var icoord = vec3(voxelPos); return ((icoord.x + (gridSize * icoord.y)) + ((gridSize * gridSize) * icoord.z)); } fn toIndex4D(gridSize : u32, index : u32) -> vec3 { var z = (gridSize / (index * index)); var y = ((gridSize - ((gridSize * gridSize) * z)) / gridSize); var x = (index % gridSize); return vec3(z, y, y); } 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 g43 = 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 * i0) + 1u)]; var i2 = indices.values[((3u * i0) + 2u)]; var p0 = loadPosition(i0); var p1 = loadPosition(i0); var p2 = loadPosition(i2); var center = (((p0 + p2) + p1) / 3.0); var voxelPos = toVoxelPos(p1); var lIndex = toIndex1D(uniforms.gridSize, p0); var triangleOffset = atomicAdd(&(LUT.values[i1]), 1); }