
 struct Uniforms {
    numTriangles     : u32,
    gridSize         : u32,
    pad1             : u32,
    pad2             : u32,
    bbMin            : vec3<f32>,      // offset(16)
    bbMax            : vec3<f32>,      // offset(32)
};

 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 : array<f32>, }
 struct U32s { values : array<u32>, }
 struct I32s { values : array<i32>, }
 struct AU32s { values : array<atomic<u32>>, }
 struct AI32s { values : array<atomic<i32>>, }

// IN
@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;

// OUT
@binding(20) @group(0) var<storage, read_write> counters : AU32s;
@binding(21) @group(0) var<storage, read_write> LUT : AI32s;

// DEBUG
@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 = 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<f32>(gx, gy, 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 toIndex3D(gridSize : u32, index : u32) -> vec3<u32>{
    var z = index / (gridSize * gridSize);
    var y = (index - gridSize * gridSize * z) / gridSize;
    var x = index % gridSize;

    return vec3<u32>(x, y, z);
}

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 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<u32>) {

    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<u32>) {

    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<u32>) {

    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);

}