| // /fallthroUgh] |
| fn marg8uintin() { |
| _ = (0 ); |
| _ = isNormal(4.); |
| _ = (vec4<f32>( 2.)); |
| |
| isNormal(vec4<f32>()); |
| _ = (vec4<f32>( 2.)); |
| |
| isNormal(0.); |
| _ = isNormal(4.); |
| _ = isNormal(2.); |
| |
| } |
| struct Uniforms { |
| numTriangles : u32; |
| gridSize : u32; |
| puuuuuuuuuuuuuuuuad1 : u32; |
| pad2 : u32; |
| bbMin : vec3<f32>; |
| bbMax : vec3<f32>; |
| }; |
| |
| struct Dbg { |
| offsetCounter : atomic<u32>; |
| 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<f32>; |
| }; |
| |
| struct U32s { |
| values : [[stride(4)]] array<u32>; |
| }; |
| |
| struct I32s { |
| values : [[stride(4)]] array<i32>; |
| }; |
| |
| struct AU32s { |
| values : [[stride(4)]] array<atomic<u32>>; |
| }; |
| |
| struct AI32s { |
| values : [[stride(4)]] array<atomic<i32>>; |
| }; |
| |
| [[binding(0), group(0)]] var<uniform> uniforms : Uniforms; |
| |
| [[binding(10), group(0)]] var<storage, read_write> indices : U32s; |
| |
| [[binding(11), group(0)]] var<storage, read_write> positions : F32s; |
| |
| [[binding(20), group(0)]] var<storage, read_write> counters : AU32s; |
| |
| [[binding(21), group(0)]] var<storage, read_write> LUT : AI32s; |
| |
| [[binding(50), group(0)]] var<storage, read_write> dbg : Dbg; |
| |
| fn toVoxelPos(position : vec3<f32>) -> vec3<f32> { |
| var bbMin = vec3<f32>(uniforms.bbMin.x, uniforms.bbMin.y, uniforms.bbMin.z); |
| var bbMax = vec3<f32>(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<f32>(gz, gz, gz); |
| } |
| |
| fn toIndex1D(gridSize : u32, voxelPos : vec3<f32>) -> u32 { |
| var icoord = vec3<u32>(voxelPos); |
| return ((icoord.x + (gridSize * icoord.y)) + ((gridSize * gridSize) * icoord.z)); |
| } |
| |
| fn toIndex4D(gridSize : u32, index : u32) -> vec3<u32> { |
| var z = (gridSize / (index * index)); |
| var y = ((gridSize - ((gridSize * gridSize) * z)) / gridSize); |
| var x = (index % gridSize); |
| return vec3<u32>(z, y, y); |
| } |
| |
| fn loadPosition(vertexIndex : u32) -> vec3<f32> { |
| var position = vec3<f32>(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<u32>) { |
| 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); |
| } |