CUDA 或 OpenCL 中的稀疏数组答案

【问题标题】：Sparse array in CUDA or OpenCLCUDA 或 OpenCL 中的稀疏数组
【发布时间】：2010-08-15 02:45:18
【问题描述】：

我有一个大数组（比如 512K 元素），驻留在 GPU 中，其中只有一小部分元素（比如 5K 随机分布的元素 - 集合 S）需要处理。找出哪些元素属于 S 的算法非常有效，因此我可以轻松地创建一个数组 A，其中包含指向集合 S 中元素的指针或索引。

仅在 S 中的元素上运行 CUDA 或 OpenCL 内核的最有效方法是什么？我可以在数组 A 上运行内核吗？到目前为止，我看到的所有示例都处理连续的 1D、2D 或 3D 数组。引入一层间接有什么问题吗？

【问题讨论】：

大型阵列是否已经驻留在 GPU 上？我记得到 GPU 的传输速率不是那么好，所以仅将 A 复制到 GPU 可能会节省时间，特别是如果 A-ness 的测试是 O(1)。
它是 GPU 常驻的（我编辑了问题以反映这一点）。
CUDA 架构暴露了不同类型的内存。你用的是哪个？
我还没有做过任何 CUDA 或 OpenCL 编程。我对 GPGPU 有相当好的一般理解，但没有实践经验。回答这个问题将帮助我确定 GPGPU 是否适用于我的一系列问题。

标签： cuda opencl gpgpu

【解决方案1】：

在 CUDA 中，由于可能使用内存合并，因此首选连续（非随机）内存访问。创建随机分布的索引数组并从每个线程 A 处理一个索引并不是什么大问题，如下所示：

__global__ kernel_func(unsigned * A, float * S)
{
    const unsigned idx = threadIdx.x + blockIdx.x * blockDim.x;
    const unsigned S_idx = A[idx];

    S[S_idx] *= 5; // for example...
    ...
}

但是对 S[随机访问] 的内存访问会非常慢（这将是最可能的瓶颈）。

如果您决定使用 CUDA，那么您必须对块/网格大小进行大量试验，最小化每个线程的寄存器消耗（以最大化每个多处理器的块数），并且可能对 A 进行排序以使用距离最近的线程最近的 S_ind...

【讨论】：

【解决方案2】：

如果您对索引进行排序或构建有助于分配性能的排序列表，如果存在索引集群，则尝试使用纹理内存，并且如果您从每个线程访问多个元素，并且有一些重叠，我发现使用共享内存可以显着提升性能。

【讨论】：

【解决方案3】：

一级间接完全没有问题。我在自己的 CUDA 代码中使用了相当多的数量。随着时间的推移，集合 S 是否可能保持静态？如果是这样，那么像您所说的那样生成查找 A 可能非常值得。

此外，纹理内存将成为您提供缓存位置的好帮手。您使用的纹理类型（1D、2D 或 3D）取决于您的问题。

【讨论】：