2

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

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

4

3 回答 3

4

在 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[random access] 的内存访问将非常慢（这将是最可能的瓶颈）。

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

于 2010-08-15T08:05:10.917 回答
1

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

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

于 2011-01-18T18:47:20.383 回答
1

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

于 2010-08-15T15:55:34.900 回答