我的CUDA程序正遭受未合并的全局内存访问.虽然第idx个线程只处理数组中的[idx] -th单元格,但是有许多间接内存访问,如下所示.
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缓存线失效
\n\nFermi 和 Kepler 架构支持两种类型的全局内存加载。完全缓存是默认模式，它尝试先命中 L1，然后命中 L2，然后命中 GMEM，加载粒度为 128 字节行。仅 L2尝试命中 L2，然后是 GMEM，加载粒度为 32 字节。对于某些随机访问模式，可以通过使 L1 无效并利用 L2 的较低粒度来提高内存效率。\xe2\x80\x93Xptxas \xe2\x80\x93dlcm=cg这可以通过使用选项进行编译来完成nvcc。
加速全局内存访问的一般准则：禁用 ECC 支持
\n\nFermi 和 Kepler GPU 支持纠错码 (ECC)，并且 ECC 默认启用。ECC 降低了峰值内存带宽，并被要求增强医学成像和大规模集群计算等应用中的数据完整性。如果不需要，可以使用 Linux 上的 nvidia-smi 实用程序（请参阅链接）或通过 Microsoft Windows 系统上的控制面板来禁用它，以提高性能。请注意，打开或关闭 ECC 需要重新启动才能生效。
\n\n在 Kepler 上加速全局内存访问的一般准则：使用只读数据缓存
\n\nKepler 具有 48KB 的数据缓存，已知仅在函数运行期间才能读取\xe2\x80\x90。使用 read\xe2\x80\x90only 路径是有益的，因为它可以卸载共享/L1 缓存路径，并且支持全速未对齐内存访问。read\xe2\x80\x90only 路径的使用可以由编译器自动管理（使用关键字const __restrict），也__ldg()可以由程序员显式管理（使用内部函数）。