concurrency - CUDA 并行过滤整数

Question

实现以下结果的规范（或实际上任何工作）方式是什么：

给定 的所有值uint32_t，我想应用一个过滤函数（谓词？）（uint32_t-> bool），以便只保留某些整数，并将它们保存到数组中。数组本身在__constant__内存中，但我想它可以移动到通用__device__内存中。接受的 s 的数量有一个实际的上限uint32_t，例如 15000。

我目前的解决方案是在代码中执行此__host__操作，其中涉及明显的for循环，但该解决方案速度慢得不切实际。相反，我想将此过滤实现为__global__内核。

我想到的解决方案是调用过滤函数并将正结果存储到数组中的内核。它写入的索引需要以某种方式原子地递增。这就是任务的关键所在——我将如何在保持性能的同时确保并发正确性？

我的代码针对 Jetson TX1，因此计算能力是5.3.

标签： concurrencycuda

Answer 1

最简单的解决方案是使用快速CUB 设备原语。更具体地说是cub::DeviceSelect. 这一项保持秩序。

或者，您可以使用您谈论的基于原子的解决方案，但请注意不会保留订单。这种解决方案在具有少量 SM 的现代 GPU 上简单且相对较快，但在高端 GPU 上可能无法很好地扩展（因为 SM 之间可能存在原子争用）。

另一种选择是使用并行扫描。这个解决方案实现起来非常复杂（因此更容易出错），但要保留 order。它可以比原子解决方案更快，并且应该可以更好地扩展（如果巧妙地编写了实现）。

最后，您可以利用阻塞和共享内存来改进您的实现（尤其是并行扫描）。阻塞可以减少 CUDA 核心之间的争用和不必要的通信。如果您不是经验丰富的 CUDA 开发人员，您可以重用 CUB 的块范围原语。