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

我的问题涉及合并全局写入到 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];
}

这里是第一个n数组的元素hash包含索引odata从最初开始更新n要点idata。显然，这会导致非常非常严重的缺乏合并。就我的代码而言，一个内核调用的哈希值与另一个内核调用的哈希值完全无关（其他内核以其他方式更新数据），因此简单地重新排序数据来优化这个特定的内核并不是一种选择。

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

纹理是一种只读机制，因此它不能直接提高 GMEM 分散写入的性能。如果你像这样“散列”：

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

（也许你的算法可以改造一下？）

那么考虑一下可能会有一些好处纹理机制 http://docs.nvidia.com/cuda/cuda-samples/index.html#simple-texture。 （您的示例本质上似乎是一维的）。

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