How to implement per voxel lighting on GPU

Question

I tried to do it like this:

1. Raytrace the scene into 2 textures, one for the position of the hit rounded to a specific resolution and a normal
2. for each pixel do lighting calculation and put information into the voxelbuffer
3. for each pixel in the position texture get the corresponding voxel from the buffer and write it to the output texture

This is what step 2 and 3 look like in code:

@group(0)
@binding(0)
var position_input: texture_storage_2d<rgba32float, read>;

@group(0)
@binding(1)
var normal_input: texture_storage_2d<rgba8sint, read>;

@group(0)
@binding(2)
var output: texture_storage_2d<rgba8unorm, write>;

@group(1)
@binding(0)
var<storage, read_write> voxels: array<Voxel>;

struct Voxel {
    position: vec3<f32>,
    color: vec3<f32>,
};

@group(1)
@binding(1)
var<storage, read_write> voxel_count: atomic<u32>;

@compute
@workgroup_size(16, 16)
fn voxelize(
    @builtin(global_invocation_id) global_id: vec3<u32>,
) {
    let output_size = vec2<u32>(textureDimensions(output));
    let position = vec2<u32>(global_id.xy);

    if (position.x >= output_size.x || position.y >= output_size.y) {
        return;
    }

    let voxel_position_full = textureLoad(position_input, position);
    // do not voxelize background
    if voxel_position_full.w == 0.0 {return;}

    let voxel_position = voxel_position_full.xyz;

    // check if voxel_position is contained in voxels
    // if not then insert it, if contained then average
    let atomic_voxel_count = atomicLoad(&voxel_count);
    for (var i = 0u; i < atomic_voxel_count; i += 1u) {
        var voxel = voxels[i];
        if all(voxel.position == voxel_position) {
            // contained, so average it
            let color = voxel_position;
            voxel.color = (voxel.color + color) * 0.5;
            return;
        }
    }
    // no voxel was found insert
    atomicAdd(&voxel_count, 1u);
    let color = abs(voxel_position + 1.0) / 10.0;
    let atomic_voxel_count_index = atomicLoad(&voxel_count);
    voxels[atomic_voxel_count_index] = Voxel(voxel_position, color);
}

@compute
@workgroup_size(16, 16)
fn summarize(
    @builtin(global_invocation_id) global_id: vec3<u32>,
    @builtin(local_invocation_id) local_id: vec3<u32>,
) {
    let output_size = vec2<u32>(textureDimensions(output));
    let position = vec2<u32>(global_id.xy);

    if (position.x >= output_size.x || position.y >= output_size.y) {
        return;
    }

    let color = vec4<f32>(0.0, 0.0, 0.0, 1.0);
    textureStore(output, position, color);

    let voxel_position_full = textureLoad(position_input, position);

    // skip background
    if voxel_position_full.w == 0.0 {return;}

    let voxel_position = voxel_position_full.xyz;
    let voxel_count = voxel_count;
    for (var i = 0u; i < voxel_count; i += 1u) {
        let voxel = voxels[i];
        if all(voxel.position == voxel_position) {
            let color = vec4<f32>(voxel.color, 1.0);
            // let color = vec4<f32>(0.0, 1.0, 0.0, 1.0);
            textureStore(output, position, color);
            return;
        }
    }
}

But this is performing very badly and does not even work right, as it has multiple memory issues like data-races.

Is there a method to fix these memory issues or is there any better method to achieve this kind of lighting?