cuda - CUDA 中间接访问导致未合并的全局内存访问

Question

我的 CUDA 程序受到未合并的全局内存访问的困扰。虽然第 idx 个线程仅处理数组中的第 [idx] 个单元，但存在许多间接内存访问，如下所示。

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

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

对于 m_fisde[idx]，我们有合并访问，但我们实际需要的是 FF[m_front[m_fside[idx]]]。有两级间接访问。

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

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

Answer 1

加速全局内存随机访问:使L1缓存线失效

费米和开普勒架构支持两种类型的全局内存加载。 完全缓存是默认模式下，它会尝试命中 L1、L2、GMEM，加载粒度为 128 字节行。 仅 L2 尝试命中 L2，然后是 GMEM，加载粒度为 32 字节。对于某些随机访问模式，可以通过使 L1 无效并利用 L2 的较低粒度来提高内存效率。这可以通过使用 –Xptxas –dlcm=cg 选项编译到 nvcc 来完成。

加速全局内存访问的一般准则:禁用 ECC 支持

Fermi 和 Kepler GPU 支持纠错码 (ECC)，并且 ECC 默认启用。 ECC 降低了峰值内存带宽，并被要求增强医学成像和大规模集群计算等应用中的数据完整性。如果不需要的话可以使用 Linux 上的 nvidia-smi 实用程序(请参阅 link )或通过 Microsoft Windows 系统上的控制面板禁用此功能可提高性能。请注意，打开或关闭 ECC 需要重新启动才能生效。

加速 Kepler 全局内存访问的一般准则:使用只读数据缓存

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