合并的全局存储器的写入使用哈希 [英] Coalesced global memory writes using hash

问题描述

我的问题是关于合并的全局写入一个集CUDA数组元素的动态变化。考虑下面的内核：

My question concerns the coalesced global writes to a dynamically changing set of elements of an array in CUDA. Consider the following kernel:

__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 的数组元素散包含的ODATA <指数/ code>来第一个 IDATA N 元素进行更新。显然，这导致可怕，可怕的缺乏聚结。在我的code的情况下，在同一内核调用散列是完全无关的散列在另一个（和其他内核更新以其他方式的数据），所以简单地重新排列数据来优化这个特定kenrel不是选项​​。

Here the first n elements of the array hash contain the indices of odata to be updated from the first n elements of idata. Obviously this leads to a terrible, terrible lack of coalescence. In the case of my code, the hash at one kernel invocation is completely unrelated to the hash at another (and other kernels update the data in other ways), so simply reordering the data to optimize this particular kenrel isn't an option.

有没有一些CUDA功能，它可以让我以改善这种状况的表现？我听到了很多谈纹理内存，但我没能翻译一下我读过成一个解决这个问题。

Is there some feature in CUDA which would allow me to improve the performance of this situation? I hear a lot of talk about texture memory, but I've not been able to translate what I've read into a solution for this problem.

推荐答案

纹理是只读的机制，所以它不能直接提高分散写入GMEM性能。如果你是散列像这样，而不是：

Texturing is a read-only mechanism, so it cannot directly improve the performance of scattered writes to GMEM. If you were "hashing" like this instead:

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

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

(perhaps your algorithm can be transformed?)

然后可能有一些利益考虑纹理机制。 （您的例子出现在自然界中1D）。

Then there might be some benefit to considering a Texture mechanism. (Your example appears to be 1D in nature).

您也可以确保您的共享内存/ L1分裂朝缓存优化。这将帮助不大，虽然有零星的写操作。

You might also make sure that your shared memory/L1 split is optimized towards cache. This won't help much with scattered writes though.

这篇关于合并的全局存储器的写入使用哈希的文章就介绍到这了，希望我们推荐的答案对大家有所帮助，也希望大家多多支持IT屋！