使用哈希技术合并全局内存写操作

Question

我的问题涉及到在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 的索引。显然，这会导致非常糟糕的内存访问效率。在我的代码中，一个内核调用中的哈希与另一个内核调用中的哈希完全不相关（其他内核以其他方式更新数据），因此仅仅通过重新排列数据来优化这个特定的内核并不能解决问题。

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

- coastal

Answer 1

纹理贴图是只读机制，因此不能直接提高对GMEM的分散写入性能。如果您改为像这样进行“哈希”：

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

也许你的算法可以改进一下？考虑使用纹理机制可能会有所帮助。（你的示例似乎是1D的）

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

- Robert Crovella

---

谢谢您的回复。事实上，我的代码中有很多地方都遇到了您所描述的情况（从全局内存中进行散乱读取）。这些读取在数组中是真正随机的（通常是稀疏的）。在这种情况下，纹理机制值得考虑吗？ - coastal

纹理贴图可能有所帮助。纹理贴图仍然取决于数据局部性和重用，以提供任何好处。如果您的分散读取的净效果是您仅读取每个位置一次，则纹理贴图无法帮助。但是，如果您可能多次读取某些位置（也许来自多个线程），则纹理贴图可能会带来一些改进。 - Robert Crovella

Answer 2

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

如果可以的话，您可以使用共享内存。您可以分配一块共享内存，并在共享内存上进行快速散写操作。然后通过合并事务将共享内存块写回全局内存。

- lashgar