用CUDA求和两个数组

Question

我正在学习CUDA,同时遵循本指南.

我还没有完成它,但我决定玩到目前为止看到的东西.

我试图重写第一个使用256个线程的例子.我想这样做,所以每个线程都在一个连续的数组切片上运行.

目标是将2个数组与1,048,576个项相加.

为了比较,这是原始代码,其中每个数组项都是根据步幅访问的:

__global__
void add(int n, float *x, float *y)
{
  int index = threadIdx.x;
  int stride = blockDim.x;
  for (int i = index; i < n; i += stride)
      y[i] = x[i] + y[i];
}

这是我的功能:

__global__
void add2(int n, float* x, float* y) {
    int sliceSize = n / blockDim.x;
    int lower = threadIdx.x * sliceSize;
    int upper = lower + sliceSize;
    for (int i = lower; i < upper; i++) {
        y[i] = x[i] + y[i];
    }
}

事实证明,最后一个片段的表现比前一个片段慢了近7倍(22ms对3ms).我认为通过在连续切片上访问它们,它会更快地执行相同的OR.

我用add<<<1, threads>>>(n, x, y)和add<<<1, threads>>>(n, x, y)(256个线程)调用函数.

值sliceSize始终是4096.在这种情况下,应该发生的是:

• threadIdx.x = 0 从0到4095
• threadIdx.x = 1 从4096到8191
• ...
• threadIdx.x = 255 从1044480到1048576

我打开了NVidia Visual Profiler,据我所知,我的内存访问模式效率不高(低全局内存负载/存储效率).第一个代码段不显示此警告.为什么会这样？

我认为第一个剪切会在数组周围跳转,从而创建一个糟糕的访问模式.实际上,它似乎很好.

我已经阅读了一些关于可视化分析器附带的内存优化的文档,但我不太明白为什么这么慢.

Answer 1

您正在探索合并内存访问和未合并内存访问之间的区别。或者我们可以简单地说“最有效”和“效率较低”的内存访问。

在 GPU 上，所有指令都在扭曲范围内执行。因此，当扭曲中的一个线程正在读取内存中的某个位置时，扭曲中的所有线程都将从内存中读取。粗略地说，最佳模式是经线中的所有线程都从相邻位置读取。这会导致 GPU 内存控制器在检查了 warp 中每个线程针对特定读取周期请求的内存地址后，可以将地址合并在一起，从而需要从缓存请求最少的行数（或从 DRAM 请求的最小段数）。

此处的幻灯片 36（或 37）以图解方式描述了这种情况。

100% 合并的情况在您的第一个代码片段中表示。从全局内存读取的示例如下：

  y[i] = x[i] + y[i];
         ^
         reading from the vector x in global memory

让我们考虑循环的第一遍，并考虑第一个扭曲的情况（即线程块中的前 32 个线程）。在这种情况下，i是由 给出的threadIdx.x。因此，线程 0 的索引为 0，线程 1 的索引为 1，依此类推。因此，每个线程正在读取全局内存中的相邻位置。假设我们错过了所有缓存，这将转换为 DRAM 读取请求，并且内存控制器可以为 DRAM 中的段（或等效地为缓存中的行）生成最小数量的请求（更准确地说：事务）。从“总线带宽利用率”为100%的意义上来说，它是最优的。在该读取周期中，请求的每个字节实际上都被扭曲中的线程使用。

“未合并”访问通常可以指不符合上述描述的任何情况。转化为上述更细粒度的“总线带宽利用率”数字，未合并访问可能有不同的程度，从略低于 100% 的最佳情况到 12.5% 或 3.125% 的最坏情况，具体取决于具体情况和图形处理器。

此处的幻灯片 44（或 45）中给出了根据本描述的最坏情况未合并访问模式示例 。这并不能准确描述最坏情况的代码片段，但对于足够大的代码片段来说，sliceSize它是等效的。代码行是相同的。考虑到相同的读取请求（for ，通过 warp 0，在循环的第一次迭代中），唯一的区别在于跨 warp 的x值：i

int sliceSize = n / blockDim.x;
int lower = threadIdx.x * sliceSize;
...
for (int i = lower; i < upper; i++) {
    y[i] = x[i] + y[i];

所以i从 开始lower，正好是threadIdx.x * sliceSize。假设sliceSize大于 1。那么第一个线程将读取位置 0。第二个线程将读取位置sliceSize。第三个线程将读取 location2*sliceSize等。这些位置之间有sliceSize距离。即使sliceSize仅为 2，该模式的效率仍然较低，因为内存控制器现在必须请求两倍数量的行或段，以满足跨 warp 0 的特定读取周期。如果sliceSize足够大，内存控制器必须请求唯一的行或每个线程的段，这是最坏情况的模式。

作为最后的说明/要点，可以对“快速分析”进行有用的观察：

• 为了在大多数情况下实现最佳合并访问，我们希望确保内存索引的计算不涉及乘以threadIdx.x1 以外的任何数量。
• 相反，如果我们可以证明threadIdx.x在任何给定的索引计算中乘以某个不等于 1 的数字，那么无论其他考虑因素如何，这几乎都表明生成的访问模式将是非最佳的。

为了清楚起见，重复一下：

index = any_constant_across_the_warp + threadIdx.x;

通常是最佳访问模式。

index = any_constant_across_the_warp + C*threadIdx.x;

通常不是最佳的访问模式。请注意，any_constant_across_the_warp可以由对数量的任意算术组成，例如：循环索引、blockIdx.?、blockDim.?和gridDim.?任何其他常量。必须考虑 2D 或 3D 线程块模式，其中threadIdx.y将被考虑在内，但将这种理解扩展到 2D 情况通常并不难。对于典型的线程块形状，为了快速分析，您通常不需要在 或 上使用恒定threadIdx.x乘数threadIdx.y。

整个讨论适用于全局内存读/写。共享内存也有最佳访问的规则，这些规则在某些方面与上面的描述相似，但在某些方面又截然不同。然而，通常情况下，全局内存的完全最佳 100% 合并模式也将是共享内存读/写的最佳模式。另一种说法是，warp 中的相邻访问通常对于共享内存也是最佳的（但这并不是共享内存唯一可能的最佳模式）。

此处已链接的演示文稿将更全面地处理该主题，网络上的许多其他演示文稿和处理也是如此。