55 lines
976 B
WebGPU Shading Language
55 lines
976 B
WebGPU Shading Language
|
|
@group(0) @binding(0)
|
||
|
|
var<storage, read_write> compare: array<u32>;
|
||
|
|
|
||
|
|
@group(0) @binding(1)
|
||
|
|
var<storage, read_write> indices: array<u32>;
|
||
|
|
|
||
|
|
@group(0) @binding(2)
|
||
|
|
var<uniform> k: u32;
|
||
|
|
|
||
|
|
@group(0) @binding(3)
|
||
|
|
var<uniform> j: u32;
|
||
|
|
|
||
|
|
@group(0) @binding(4)
|
||
|
|
var<uniform> totalCount: u32;
|
||
|
|
|
||
|
|
@compute @workgroup_size(256)
|
||
|
|
fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
|
||
|
|
let idx = global_id.x;
|
||
|
|
|
||
|
|
let ixj = idx ^ j;
|
||
|
|
|
||
|
|
if (idx >= totalCount || ixj <= idx || ixj >= totalCount) {
|
||
|
|
return;
|
||
|
|
}
|
||
|
|
if (ixj > idx) {
|
||
|
|
let ascending = (idx & k) == 0u;
|
||
|
|
|
||
|
|
let dist_idx = compare[idx];
|
||
|
|
let dist_ixj = compare[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 = compare[idx];
|
||
|
|
let tempIndex = indices[idx];
|
||
|
|
|
||
|
|
compare[idx] = compare[ixj];
|
||
|
|
compare[ixj] = tempDist;
|
||
|
|
|
||
|
|
indices[idx] = indices[ixj];
|
||
|
|
indices[ixj] = tempIndex;
|
||
|
|
}
|
||
|
|
}
|
||
|
|
}
|