import wgpu_3dgs_core::{
gaussian::Gaussian,
gaussian_transform::GaussianTransform,
model_transform::ModelTransform,
};
import package::selection::consts;
@group(0) @binding(0)
var<uniform> op: u32;
@group(0) @binding(1)
var<storage, read> source: array<u32>;
@group(0) @binding(2)
var<storage, read_write> dest: array<atomic<u32>>;
@group(0) @binding(3)
var<uniform> model_transform: ModelTransform;
@group(0) @binding(4)
var<uniform> gaussian_transform: GaussianTransform;
@group(0) @binding(5)
var<storage, read> gaussians: array<Gaussian>;
override workgroup_size: u32;
@compute @workgroup_size(workgroup_size)
fn main(@builtin(global_invocation_id) id: vec3<u32>) {
let index = id.x;
if (index >= arrayLength(&dest)) {
return;
}
if op == consts::op_union {
atomicOr(&dest[index], source[index]);
} else if op == consts::op_intersection {
atomicAnd(&dest[index], source[index]);
} else if op == consts::op_symmetric_difference {
atomicXor(&dest[index], source[index]);
} else if op == consts::op_difference {
atomicStore(&dest[index], source[index] & ~atomicLoad(&dest[index]));
} else if op == consts::op_complement {
atomicStore(&dest[index], ~atomicLoad(&dest[index]));
}
}