2

我的内核使用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：目标架构是费米和开普勒。

4

1 回答 1

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 实用程序（请参阅链接）或通过 Microsoft Windows 系统上的控制面板禁用它以提高性能。请注意，打开或关闭 ECC 需要重新启动才能生效。

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

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

于 2013-03-02T07:18:51.927 回答