我的 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 中找到一些数据模式，以使其成为直接顺序访问，但发现它们几乎是“随机的”。
有没有办法解决这个问题？
加速全局内存随机访问：使 L1 缓存行无效
Fermi 和 Kepler 架构支持来自全局内存的两种类型的负载。完全缓存是默认模式，它尝试在 L1、L2、GMEM 中命中，加载粒度为 128 字节行。L2-only尝试在 L2 中命中，然后是 GMEM，加载粒度为 32 字节。对于某些随机访问模式，可以通过使 L1 无效并利用 L2 的较低粒度来提高内存效率。这可以通过编译–Xptxas –dlcm=cg选项来完成nvcc。
加速全局内存访问的一般准则：禁用 ECC 支持
Fermi 和 Kepler GPU 支持纠错码 (ECC)，默认启用 ECC。ECC 降低了峰值内存带宽，并被要求增强医学成像和大规模集群计算等应用程序中的数据完整性。如果不需要，可以使用 Linux 上的 nvidia-smi 实用程序（请参阅链接）或通过 Microsoft Windows 系统上的控制面板禁用它以提高性能。请注意，打开或关闭 ECC 需要重新启动才能生效。
在 Kepler 上加速全局内存访问的一般准则：使用只读数据缓存
Kepler 具有一个 48KB 的缓存，用于存储已知在函数执行期间为只读的数据。使用只读路径是有益的，因为它减轻了共享/L1 缓存路径的负担，并且支持全速非对齐内存访问。只读路径的使用可以由编译器自动管理（使用const __restrict关键字）或__ldg()由程序员显式管理（使用内在函数）。