WebGPU 基数键值排序

本包实现了基数排序的GPU版本。有关通用基数排序的良好介绍，请参阅此处：http://www.codercorner.com/RadixSortRevisited.htm

此处实现的GPU基数排序是对fuchsia仓库中找到的Vulkan基数排序的重实现：https://fuchsia.googlesource.com/fuchsia/+/refs/heads/main/src/graphics/lib/compute/radix_sort/。

目前仅实现了32位键值对的排序。它可以用来排序无符号整数和正浮点数。有关更多详细信息，请参阅限制。键值按升序排序。

最初是为我们的3D高斯Splatting渲染器实现的，用于实时根据Splat的深度进行排序。您可以在这个网页演示中看到它的实际应用。

示例

// find best subgroup size
let subgroup_size = guess_workgroup_size(&device, &queue).await.unwrap();
let sorter = GPUSorter::new(&device, subgroup_size);

// setup buffers to sort 100 key-value pairs
let n = 100;
let sort_buffers = sorter.create_sort_buffers(&device, NonZeroU32::new(n).unwrap());

let keys_scrambled: Vec<u32> = (0..n).rev().collect();
let values_scrambled:Vec<u32> = keys_scrambled.clone();

let mut encoder = device.create_command_encoder(&wgpu::CommandEncoderDescriptor {label: None});

upload_to_buffer(
    &mut encoder,
    &sort_buffers.keys(),
    &device,
    keys_scrambled.as_slice(),
);
upload_to_buffer(
    &mut encoder,
    &sort_buffers.values(),
    &device,
    values_scrambled.as_slice(),
);

// sorter.sort(&mut encoder, &sort_buffers);
sorter.sort(&mut encoder,&queue,&sort_buffers,None);
queue.submit([encoder.finish()]);

// key and value buffer is now sorted.

还支持间接调度。请参阅examples/sort_indirect.rs；

基准测试

为了测量性能，我们进行了1000次键值对的排序，并报告了每次运行的平均持续时间。对不同数量的对进行了测量。有关更多详细信息，请参阅benches/sort.rs。

限制

本排序器存在一些限制，以下将进行解释。

子组

此渲染器使用子群来减少同步并提高性能。遗憾的是，目前WebGPU/wgpu不支持子群操作。

为了克服这个问题，我们通过尝试不同的子群“猜测”子群大小，并选择最大的可行子群（参见utils::guess_workgroup_size）。这在几乎所有情况下都有效，但可能会失败，因为子群大小可能会随时间变化。一旦支持子群，这个问题将会得到解决。状态可以在这里找到：这里。

浮点数

排序算法将值解释为整数，并按升序对键进行排序。非负浮点值可以解释为无符号整数，而不会影响排序顺序。因此，此排序器可以用于排序32位浮点键。请注意，NaN和Inf值会导致意外结果，因为这些值也被解释为整数。排序浮点值的示例可以在这里找到。