由于CUDA中的间接访问引起的未合并全局内存访问

Question

由于CUDA中的间接访问引起的未合并全局内存访问

6

我的CUDA程序遭受了不连续的全局内存访问问题。虽然第idx个线程只处理数组中的[idx]单元格，但下面显示了许多间接内存访问。

int idx=blockDim.x*blockIdx.x+threadIdx.x;

.... = FF[m_front[m_fside[idx]]];

对于m_fside[idx]，我们已经合并了访问，但实际上我们需要的是FF[m_front[m_fside[idx]]]。存在两级间接访问。

我试图找到m_front或m_fside中数据的一些模式，以使其成为直接顺序访问，但发现它们几乎是“随机”的。

有没有可能解决这个问题？

- thierry

1

这实际上与稀疏矩阵寻址问题相同，对于如何改进此问题已经进行了相当多的工作。您可以从查看有关GPUS上稀疏矩阵操作的文献中获得一些想法。 - talonmies

如果访问中存在任何局部性，则可能会对此问题感兴趣。 - Robert Crovella

1

@RobertCrovella... 上面链接的答案中提供的纹理机制链接已经失效。您能否更新一下链接？ - sgarizvi

1

@sgar91 链接已经修复。 - Robert Crovella

1

@thierry 请查看我的修订和改进的答案。 - Vitality

@RobertCrovella 纹理当然是一个好的解决方案。但是我的数据很大，数组也很多。我担心如果利用纹理内存，我将会有很多纹理内存替换操作，这可能会带来新的开销。 - thierry

1个回答

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- Vitality · Accepted Answer

加速全局内存随机访问：使L1缓存行无效

Fermi和Kepler架构支持从全局内存中加载两种类型的数据。默认模式是“完全缓存”，它会尝试在L1、L2、GMEM中查找，加载粒度为128字节行。而“仅L2”模式则会尝试在L2、GMEM中查找，加载粒度为32字节。对于某些随机访问模式，可以通过使L1失效并利用L2更低的粒度来提高内存效率。这可以通过使用nvcc编译时的–Xptxas –dlcm=cg选项来实现。

加速全局内存访问的一般指南：禁用ECC支持

Fermi和Kepler GPU支持纠错码（ECC），默认情况下启用ECC。ECC可以降低峰值内存带宽，并在需要提高数据完整性的应用程序中使用，例如医学成像和大规模集群计算。如果不需要，可以使用Linux上的nvidia-smi实用程序（请参见link）或Microsoft Windows系统上的控制面板禁用它以提高性能。请注意，切换ECC开关需要重新启动才能生效。

Kepler具有48KB缓存，用于存储在函数期间已知为只读的数据。使用只读路径是有益的，因为它可以卸载共享/L1缓存路径并且支持全速非对齐内存访问。编译器可以自动管理只读路径（使用const __restrict关键字），程序员也可以显式管理（使用__ldg()内置函数）。