3

我的问题涉及合并的全局写入到 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的索引。显然，这导致了可怕的、可怕的缺乏结合。在我的代码的情况下，一个内核调用的哈希与另一个内核调用的哈希完全无关（其他内核以其他方式更新数据），因此简单地重新排序数据以优化这个特定的 kenrel 不是一种选择。odatanidata

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

4

2 回答 2

3

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

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

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

那么考虑纹理机制可能会有一些好处。（您的示例本质上似乎是一维的）。

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

于 2012-10-17T16:16:23.510 回答
0

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

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

于 2012-10-18T16:28:58.180 回答