struct VoxelStructureNode
{
    pointers: array<u32, 64>
}

struct VoxelColorNode
{
    pointers: array<vec4<f32>, 64>
}

struct CacheNodeRequest
{
    location_info: u32,
    request_count: u32
}

struct CacheNodeUsageBuffer
{
    touch_time: u32,
    parent: u32,
}

// Rendering kernels
struct ChunkInfo
{
    mvp: mat4x4<f32>,
    eye_pos: vec3<f32>,
    root_color: vec4<f32>,
    root_subdiv: u32,
    frame_timestamp: u32,
}

struct StructureNode
{
    pointers: array<u32, 64>
}

struct RequestElement
{
    pointers: array<atomic<u32>, 64>
}

struct ColorNode
{
    colors: array<vec4<f32>, 64>
}

struct VertexOutput
{
    @builtin(position) postion: vec4<f32>,
    @location(0) world_loc: vec3<f32>,
    @location(1) cam_pos: vec3<f32>,
}

var<immediate> constants: ChunkInfo;
//var<push_constant> constants: ChunkInfo;

@vertex
fn chunk(@builtin(vertex_index) index: u32) -> VertexOutput
{
    let cube_vertices = array<vec3<f32>, 8>(
        vec3<f32>(0., 0., 0.),
        vec3<f32>(0., 0., 1.),
        vec3<f32>(1., 0., 1.),
        vec3<f32>(1., 0., 0.),

        vec3<f32>(0., 1., 0.),
        vec3<f32>(0., 1., 1.),
        vec3<f32>(1., 1., 1.),
        vec3<f32>(1., 1., 0.),
    );

    let cube_faces = array<u32, 24>(
        // Bottom face
        1, 0, 2, 3,

        // Top face
        4, 5, 7, 6,

        // Side faces
        0, 1, 4, 5,
        1, 2, 5, 6,
        2, 3, 6, 7,
        3, 0, 7, 4,
    );

    let quad_index = index / (3 * 2);
    let triangle_index = index % (3 * 2);
    let triangle_map = array<u32, 6>(
        0, 1, 2, 1, 3, 2
    );

    let vertex = cube_vertices[cube_faces[quad_index * 4 + triangle_map[triangle_index]]];
    let output_vertex = constants.mvp * vec4<f32>(vertex, 1.0f);

    var output: VertexOutput;
    output.postion = output_vertex;
    output.world_loc = vertex;
    output.cam_pos = constants.eye_pos;

    //let output = vec4<f32>(vertex, 1.0f);
    return output;
}

@group(0) @binding(0) var<storage, read> structure_buffer: array<StructureNode>;
@group(0) @binding(1) var<storage, read> color_buffer: array<ColorNode>;
@group(0) @binding(2) var<storage, read_write> request_buffer: array<RequestElement>;
@group(0) @binding(3) var<storage, read_write> usage_buffer: array<u32>;


fn box_inter(pos: vec3<f32>, ray_dir: vec3<f32>, box_min: vec3<f32>, box_max: vec3<f32>) -> vec2<f32>
{
    let box_min_t = (box_min - pos) / ray_dir;
    let box_max_t = (box_max - pos) / ray_dir;

    let near_ts = min(box_min_t, box_max_t);
    let far_ts = max(box_min_t, box_max_t);

    let far_t = min(min(far_ts.x, far_ts.y), far_ts.z);
    let near_t = max(max(near_ts.x, near_ts.y), near_ts.z);

    return vec2(near_t, far_t);
}

fn sdf(voxel: vec3<i32>) -> bool
{
    let len = length(vec3<f32>(voxel) - vec3(128)) / 128.;
    return len <= 1.;
}

fn min_vec(x: vec3<f32>) -> f32
{
    return min(x.x, min(x.y, x.z));
}

fn min_mask(x: vec3<f32>) -> vec3<bool>
{
    let min = min(x.x, min(x.y, x.z));

    return vec3<bool>(min == x.x, min == x.y, min == x.z);
}

fn node_subdivided(node: u32) -> bool
{
    return (node >> 31) != 0;
}

fn node_pointer_valid(node: u32) -> bool
{
    return ((node >> 30) & 1) != 0;
}

fn node_pointer(node: u32) -> u32
{
    return node & 0x3FFFFFFF;
}

fn voxel_from_wall(position: vec3<f32>, ray_dir: vec3<f32>) -> vec3<i32>
{
    let integers = round(position);
    let wall_mask = min_mask(abs(position - vec3<f32>(integers)));
    let offsets = select(vec3<f32>(-0.5), vec3<f32>(0.5), ray_dir > vec3(0.));
    return vec3<i32>(floor(position + select(vec3<f32>(0.), offsets, wall_mask)));
}

fn traverse(ray_dir: vec3<f32>, ray_origin: vec3<f32>, root_color: vec4<f32>, root_subdiv: bool) -> vec4<f32>
{
    // Simple FVT
    let t_off = abs(1. / ray_dir);

    // Start location
    let voxel_dir = select(vec3(-1), vec3(1), ray_dir >= vec3(0.));
    var pos_origin = clamp(ray_origin * 256., vec3(0.), vec3(256. - 1.));
    var voxel = vec3<i32>(pos_origin);
    var last_voxel = voxel;

    let wall_offset = select(vec3(0), vec3(1), ray_dir > vec3(0.));

    var dfs_stack = array<u32, 5>(0, 0, 0, 0, 0);

        // Current depth of the node we are exploring
        var current_depth = 0;

        // Index of the current node's data
        var current_node = u32(0);

        // Current node size
        var node_size = 4 * 4 * 4 * 4; // 128

        // Size of a child of this node
        var child_size = node_size / 4;

        // Lut of the node_size per depth
        var node_size_lut = array<i32, 5>(
                4 * 4 * 4 * 4,
                4 * 4 * 4,
                4 * 4,
                4,
                1,
                );

    let depth_limit = 1;
    for(var iter = 0; iter < 256; iter ++)
    {

        // Our ray is currently touching a voxel.
        // Descend to the lowest node that contains this voxel

        // Position of the child we are in
        var child_pos = (vec3<u32>(voxel) >> vec3<u32>((4 - u32(current_depth + 1)) * 2)) & vec3<u32>(3); // Hardcode for 4-tree
        var child_index = child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4;

        // Current node has been used. report
        usage_buffer[current_node] = constants.frame_timestamp;
        
        while(
            node_subdivided(structure_buffer[current_node].pointers[child_index]) 
            && current_depth < depth_limit)
        {
            if(!node_pointer_valid(structure_buffer[current_node].pointers[child_index]))
            {
                atomicAdd(&request_buffer[current_node].pointers[child_index], 1);
                break;
            }
            // Child node is subdivided, we go in, save position in stack
            current_node = node_pointer(structure_buffer[current_node].pointers[child_index]);
            usage_buffer[current_node] = constants.frame_timestamp;
            current_depth += 1;
            dfs_stack[current_depth] = current_node;
            node_size = node_size_lut[current_depth];

            child_pos = (vec3<u32>(voxel) >> vec3<u32>((4 - u32(current_depth + 1)) * 2)) & vec3<u32>(3); // Hardcode for 4-tree
            child_index = child_pos.x + child_pos.y * 4 + child_pos.z * 4 * 4;
            child_size = node_size / 4;
        }


        // At this point current_depth is the depth of the node that contains the voxel
        // child_pos and child_index relate to the specific child of the node that contains this voxel
        // It is guaranteed that the child is leave

        // Check current leave's color
        let color = color_buffer[current_node].colors[child_index];
        if(color.w != 0.) // Not transparent
        {
            let k = child_pos.x + child_pos.y + child_pos.z;
            let x = select(0.5, 1., k % 2 == 0);
            return x * color / f32(current_depth);
        }

        // Voxel and whole child containing it is empty
        
        // Perform a step through the children of the node
        let child_position = (voxel / child_size) * child_size;
        let far_corner = child_position + wall_offset * child_size;
        let far_ts = (vec3<f32>(far_corner) - pos_origin) / ray_dir; // TODO: Turn into fma
        let far_t = min(min(far_ts.x, far_ts.y), far_ts.z);

        let next_child_min = select(child_position, child_position + wall_offset * child_size, vec3(far_t) == far_ts);
        let next_child_max = next_child_min + vec3(child_size);
        
        // The ray (far_t) is now touching the new child to explore
        // Find out which actual voxel we are touching
        let previous_voxel = voxel;
        let float_voxel = clamp(pos_origin + far_t * ray_dir, vec3<f32>(next_child_min), vec3<f32>(next_child_max));
        /*
        voxel = vec3<i32>(
            floor(
            select(
                float_voxel - vec3(0.5),
                float_voxel + vec3(0.5),
                ray_dir > vec3(0.)
            ))
        );
        */
        //voxel = vec3<i32>(round(float_voxel));
        voxel = voxel_from_wall(float_voxel, ray_dir);
        if(any(voxel < vec3(0)) || any(voxel >= vec3(256)))
        {
            //return vec4(f32(iter) / 100.);
            discard;
        }

        // We touched a voxel as if we explored blocks sized by the child size of the current node.
        // But we might have exited the current node.
        
        // If this is the case we have to walk back up the tree
        // And then back down to the next node over

        // As such we find the lowest ancestor that can contain both the privous voxel (in node) and the new voxel (out of node)
        let bit_diffs = voxel ^ previous_voxel;
        let bit_diffs_lowest = bit_diffs.x | bit_diffs.y | bit_diffs.z;

        let flb = ((countLeadingZeros(bit_diffs_lowest) - 24) / 2);
        let common_depth = flb;

        current_depth = common_depth;
        node_size = node_size_lut[current_depth];
        child_size = node_size / 4;
        current_node = dfs_stack[current_depth];
        
        // Figure out current voxel position
        //voxel = vec3<i32>(ray_origin + ray_dir * t); 

    }
    return vec4<f32>(1., 0., 1., 1.);

}

@fragment
fn fragment(in: VertexOutput) -> @location(0) vec4<f32>
{
    let ray_dir = normalize(in.world_loc - in.cam_pos);
    let interp = box_inter(in.cam_pos, ray_dir, vec3(0.), vec3(1));
    let ray_origin = in.cam_pos + ray_dir * (max(0., interp.x));

    return traverse(ray_dir, ray_origin, constants.root_color, constants.root_subdiv != 0);
    return vec4(interp.y / 10.);
}

// Cache managment kernels

// Request buffer managment
// Compaction
@workgroup_size(64)
fn main(
    @builtin(global_invocation_id) gid: vec3<u32>
)
{
    let cache_size = 1024;
    let index = gid.x;

    for()
}