CUDA 对全局内存的低效访问模式答案

【问题标题】：CUDA inefficient access pattern to global memoryCUDA 对全局内存的低效访问模式
【发布时间】：2018-05-04 16:55:03
【问题描述】：

我编写了一个 CUDA 内核，它执行两个数组 arr1 和 arr2 的数组加法。 arr1的哪个索引应该加上arr2的哪个索引存储在数组idx中的信息。

这是一个代码示例：

__global__ add(float* arr1, float* arr2, int* idx, int length)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    // each thread performs (length) additions,
    // arr2 is (lenght) times larger than arr1
    for (int j=threadIdx.x; j<length*blockDim.x; j+=blockDim.x)
    {
        arr1[i] += arr2[ idx[blockIdx.x*blockDim.x + j] ]; // edited here
    }
}

该代码产生正确的输出，但几乎不比具有 8 个线程的 CPU 上的 openmp 并行代码快。我尝试了不同的块大小。

我怀疑对arr2 的访问模式效率低下，因为arr2 在全局内存中并且是准随机访问的——数组idx 包含唯一的、已排序但不连续的索引（可能是2, 3, 57, 103, ...)。因此，没有利用 L1 缓存。此外，数组非常大，无法完全放入共享内存中。

有没有办法绕过这个障碍？你知道如何优化arr2的访问模式吗？

【问题讨论】：

是什么让您认为访问arr2 是瓶颈？对idx 的完全未合并的访问呢？
好点。尽管如此，一个线程至少访问idx 的length 连续元素，而arr2 的访问模式中没有任何顺序。
我认为您需要阅读有关内存合并的内容
是的，谢谢，我做了一些研究并更新了问题。现在应该合并对idx 的访问，并且我的代码运行速度是原来的两倍。现在我很确定访问arr2 是瓶颈。
编辑后的代码不可能产生与您问题中的原始代码相同的结果

标签： caching cuda

【解决方案1】：

您可以在这里尝试做几件事：

全局内存很慢，您可以尝试将数组加载到每个块的共享内存中。因此，例如，您将为每个块加载部分数组。
如果有任何东西被定义为常量，您应该将它移到常量内存中。
尝试展开循环。
一定要尝试使用不同的流参数以及不同的块/线程组合启动多个内核。

希望这能让您对优化方式有所了解:)

【讨论】：

非常感谢！连续将 arr2 的块复制到共享内存并对其进行操作确实显着加快了执行时间。
您也可以尝试使用#pragma unroll 展开循环，这应该会加快一点速度。划分内核以在不同的流中同时执行内核也会加快速度，但会花费您宝贵的开发时间:)
我的应用程序的瓶颈现在在其他地方，cuda 内核现在已经足够快了。但只是出于兴趣，我会尝试你的建议；）
在 nsight 上使用 Profile Cuda 应用程序 -> 在 Visual Studio 中开始性能分析。运行所有的 cuda 实验，这将帮助您找出应用程序的瓶颈。祝你好运！