从 CUDA 内核随机访问小型常量数组的工作原理

Question

我的内核使用 float 大小为 8 x 8 的数组，下面是随机访问模式。

// inds - big array of indices in range 0,...,7
// flts - 8 by 8 array of floats

// kernel essentially processes large 2D array by looping through slow coordinate
// and having block/thread parallelization of fast coordinate.

__global__ void kernel (int const* inds, float const* flt, ...)
{
    int idx = threadIdx.x + blockDim.x * blockIdx.x;  // Global fast coordinate
    int idy;                                          // Global slow coordinate
    int sx = gridDim.x * blockDim.x;                  // Stride of fast coordinate

    for ( idy = 0; idy < 10000; idy++ )       // Slow coordinate loop
    {
        int id = idx + idy * sx;              // Global coordinate in 2D array

        int ind = inds[id];                   // Index of random access to small array

        float f0 = flt[ind * 8 + 0];
        ...
        float f7 = flt[ind * 8 + 7];

        NEXT I HAVE SOME ALGEBRAIC FORMULAS THAT USE f0, ..., f7
    }
}

访问flt 数组的最佳方式是什么？

1. 不要通过flt，使用__const__内存。我不确定当不同线程访问不同数据时 const 内存有多快。
2. 如上使用。不会使用负载统一，因为线程访问不同的数据。由于缓存，它会很快吗？
3. 复制到共享内存并使用共享内存数组。
4. 使用纹理。从未使用过纹理...这种方法能很快吗？

对于共享内存，转置flt 数组可能会更好，即以这种方式访问​​它以避免银行冲突：

float fj = flt_shared[j * 8 + ind]; // where j = 0, ..., 7

PS：目标架构是 Fermi 和 Kepler。

Answer 1

“最佳”方式还取决于您正在处理的架构。我个人在 Fermi 和 Kepler 上随机访问的经验（由于使用映射 inds[id]，您的访问似乎有点随机）是 L1 现在非常快，在许多情况下最好继续使用全局内存而不是共享内存或纹理内存。

加速全局内存随机访问：使 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 实用程序（请参阅link）或通过 Microsoft Windows 系统上的控制面板禁用以提高性能。请注意，打开或关闭 ECC 需要重新启动才能生效。

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

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