使用哈希合并全局内存写入答案

【问题标题】：Coalesced global memory writes using hash使用哈希合并全局内存写入
【发布时间】：2012-10-17 15:53:38
【问题描述】：

我的问题涉及合并全局写入到 CUDA 中数组的一组动态变化的元素。考虑以下内核：

__global__ void
kernel (int n, int *odata, int *idata, int *hash)
{
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n)
    odata[hash[i]] = idata[i];
}

这里数组hash 的第一个n 元素包含要从idata 的第一个n 元素更新的odata 的索引。显然，这导致了可怕的、可怕的缺乏结合。在我的代码的情况下，一个内核调用的哈希与另一个内核调用的哈希完全无关（其他内核以其他方式更新数据），所以简单地重新排序数据以优化这个特定的 kenrel 不是一种选择。

CUDA 中是否有一些功能可以让我改善这种情况的性能？我听到很多关于纹理内存的讨论，但我无法将我读到的内容转化为解决这个问题的方法。

【问题讨论】：

标签： c++ c cuda gpgpu

【解决方案1】：

Texturing 是一种只读机制，因此不能直接提高对 GMEM 的分散写入的性能。如果您改为这样“散列”：

odata[i] = idata[hash[i]];

（也许你的算法可以转换？）

那么考虑Texture mechanism 可能会有一些好处。（您的示例在本质上似乎是一维的）。

您还可以确保您的共享内存/L1 拆分针对缓存进行了优化。不过，这对分散的写入没有多大帮助。

【讨论】：

感谢您的回复。实际上，我的代码中有很多地方都遇到了您所描述的情况（来自全局内存的分散读取）。这些读取在数组中是真正随机的（通常是稀疏的）。在这种情况下，纹理机制值得研究吗？
纹理可能在这种情况下有所帮助。纹理化仍然依赖于数据局部性和重用以提供任何好处。如果您的分散读取的净效果是您只读取每个位置一次，那么纹理将无济于事。但是，如果您可能不止一次地读取某些位置（可能来自多个线程），那么纹理可能会带来一些改进。

【解决方案2】：

你能限制哈希结果的范围吗？例如，您可能知道线程的前 1K 次迭代只能访问 odata 的 0 到 8K 范围内。

如果是这样，您可以使用共享内存。您可以分配共享内存块并临时在共享内存上进行快速分散写入。然后在合并事务中将共享内存块写回全局内存。

【讨论】：