如何在OpenCL/CUDA中实现稀疏数组的高效压缩？

当然有高效的GPU Kernel方案来处理这个稀疏数组压缩问题！你提到归约算法不适用完全正确——我们需要的是**并行扫描（前缀和）**这类算法，这正是GPU处理数据重排类任务的拿手好戏，比CPU串行方案的效率提升在大数据量下会非常显著。

核心思路：并行扫描+数据重排 #

这类问题的本质是给每个非零元素确定它在压缩后数组中的目标位置，然后完成数据迁移。具体分三步：

1. 生成掩码数组
   每个线程并行检查对应位置的元素是否非零，生成一个由0和1组成的掩码数组：原元素非零则掩码为1，否则为0。比如你的示例数组对应的掩码是[1,0,0,0,1,0,0,1,0,0,0,0,1,0,0,0,1,0,1,0,0,1,0,0,0,0,1]。

2. 计算独占前缀和
   对掩码数组执行独占前缀和（Exclusive Scan），这样每个非零元素对应的前缀和结果，就是它在压缩后数组中的目标索引。比如示例中：

   • 第一个非零元素（索引0）的前缀和是0，对应压缩数组的第0位
   • 第二个非零元素（索引4）的前缀和是1，对应压缩数组的第1位
   • 以此类推，所有非零元素都会得到唯一且连续的目标位置。

3. 并行数据重排
   每个线程检查自身负责的元素是否非零，如果是，就将其复制到前缀和给出的目标位置。最后可以把压缩数组的剩余位置清零（如果需要的话）。

为什么这比CPU串行方案高效？ #

你提到CPU的串行O(N)方案开销大，这完全合理——CPU只能单线程（或少量线程）遍历数组，而GPU可以同时启动数千个线程并行处理。虽然理论时间复杂度都是O(N)，但GPU的并行执行能把实际运行时间降到CPU的几十分之一甚至更低，尤其当数组规模达到百万、亿级元素时。

代码实现示例 #

如果你用CUDA生态，最便捷的方式是用Thrust库（NVIDIA官方的并行算法库），它已经封装好了优化的实现：

#include <thrust/device_vector.h>
#include <thrust/copy.h>
#include <thrust/functional.h>

int main() {
    // 初始化设备端输入数组（假设已经从主机端拷贝过来）
    thrust::device_vector<int> d_input = {9, 0, 0, 0, 7, 0, 0, 3, 0, 0, 0, 0, 5, 0, 0, 0, 8, 0, 2, 0, 0, 4, 0, 0, 0, 0, 5};
    thrust::device_vector<int> d_output(d_input.size());

    // 把非零元素复制到输出数组的开头
    auto non_zero_end = thrust::copy_if(
        d_input.begin(), d_input.end(),
        d_output.begin(),
        thrust::not_equal_to<int>(0)
    );

    // 可选：将输出数组剩余位置清零
    thrust::fill(non_zero_end, d_output.end(), 0);

    // 此时d_output就是你需要的压缩后的数组
    return 0;
}

如果需要手写自定义Kernel（比如特殊硬件适配需求），可以分三步实现：

1. 掩码生成Kernel：

__global__ void generate_mask(int* input, int* mask, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        mask[idx] = (input[idx] != 0) ? 1 : 0;
    }
}

2. 前缀和计算（可以用Thrust的thrust::exclusive_scan，或自己实现并行扫描Kernel，比如Hillis-Steele算法）
3. 数据重排Kernel：

__global__ void scatter_nonzeros(int* input, int* output, int* prefix_sum, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n && input[idx] != 0) {
        output[prefix_sum[idx]] = input[idx];
    }
}

补充说明 #

• 关于时间复杂度：这类数据重排问题理论上不可能低于O(N)——每个元素至少需要被访问一次，但GPU的并行O(N)和CPU的串行O(N)在实际性能上天差地别。
• 优化点：如果非零元素比例极低，可以考虑先筛选出非零元素的索引再做重排，不过对于25%的比例，扫描+重排的方案已经足够高效。

内容的提问来源于stack exchange，提问作者MaiaVictor