联合的全球记忆使用哈希写作

Coalesced global memory writes using hash

本文关键字：哈希写 记忆      更新时间：2023-10-16

我的问题涉及融合的全球写入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元素包含odata的索引，该索引将从idata的第一个n元素中更新。显然，这导致了可怕的合并不足。对于我的代码，一个内核调用的哈希人完全与另一个哈希完全无关（其他内核以其他方式更新数据），因此，只需重新排序数据以优化此特定的Kenrel，这不是一个选项。<<<<<<

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

纹理是一种仅读取机制，因此它不能直接提高散射的写入对gmem的性能。如果您是这样的"哈希"：

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

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

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

您还可以确保将共享内存/L1拆分优化针对缓存。

，这对分散的写入无济于事。

您可以限制哈希结果的范围吗？例如，您可能知道，线程的第一个1K迭代仅在odata的0到8K范围内访问。

如果是这样，您可以使用共享内存。您可以分配共享内存的块，并暂时在共享内存上写入快速散布。然后将共享内存块写回合并交易中的全局内存。

相关文章：