@group(0) @binding(0)
var<storage, read_write> gridHashes: array<u32>;

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

@group(0) @binding(2)
var<storage, read_write> kArray: array<u32>;

@group(0) @binding(3)
var<storage, read_write> jArray: array<u32>;

@group(0) @binding(4)
var<storage, read_write> indices: array<u32>;

@group(0) @binding(5)
var<uniform> totalCount: u32;

@compute @workgroup_size(256)
fn main(@builtin(global_invocation_id) global_id : vec3<u32>) {

	let idx = global_id.x;

	let threadPassIndex = threadPassIndices[idx];

	threadPassIndices[idx] = threadPassIndices[idx] + 1u;

	let j = jArray[threadPassIndex];

	let k = kArray[threadPassIndex];

	let ixj = idx ^ j;

	if (ixj <= idx || ixj >= totalCount) {
		return;
	}

	if (ixj > idx) {
		let ascending = (idx & k) == 0u;

		let dist_idx = gridHashes[idx];
		let dist_ixj = gridHashes[ixj];

		var swap = false;

		if (ascending) {
			if (dist_idx > dist_ixj) {
				swap = true;
			}
		} else {
			if (dist_idx < dist_ixj) {
				swap = true;
			}
		}

		if (swap) {
			let tempDist = gridHashes[idx];
			let tempIndex = indices[idx];

			gridHashes[idx] = gridHashes[ixj];
			gridHashes[ixj] = tempDist;

			indices[idx] = indices[ixj];
			indices[ixj] = tempIndex;
		}
	}
}