用CUDA求两个数组的和

Summing two arrays with CUDA

本文关键字：两个 数组 CUDA      更新时间：2023-10-16

我正在学习CUDA，同时遵循本指南。

我还没有完成，但我决定用我目前所看到的来打一点。

我试图重写使用256个线程的第一个示例。我想让它让每个线程都在数组的一个连续切片上操作。

目标是将2个数组与1048576个项目相加。

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

__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倍(22毫秒对3毫秒)。我认为，通过在一个连续的切片上访问它们，它可以更快地执行相同的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，我了解到我的内存访问模式效率不高(全局内存加载/存储效率低)。此警告不存在于第一个代码段中。为什么会出现这种情况？

我认为第一个剪切会在数组周围跳跃，从而创建一个糟糕的访问模式。事实上，这似乎很好。

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

您正在探索合并内存访问和非恢复内存访问之间的区别。或者我们可以简单地说"最高效"answers"不太高效"的内存访问。

在GPU上，所有指令都在曲速范围内执行。因此，当一个扭曲中的一个线程正在读取内存中的某个位置时，扭曲中的所有线程都在从内存中读取。粗略地说，最佳模式是当经线中的所有线都从相邻位置读取时。这导致这样一种情况，即GPU存储器控制器在针对特定读取周期检查由扭曲中的每个线程请求的存储器地址之后，可以将地址合并在一起，以导致需要从高速缓存请求的最小行数(或从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%的意义上讲，它是最优的。请求的每个字节实际上都被warp中的线程在读取周期中使用。

"未恢复"访问通常指不符合上述描述的任何情况。转化为上述细粒度的"总线带宽利用率"数字，根据具体情况和GPU，非恢复访问可能具有不同程度，从略低于100%的最佳情况到12.5%或3.125%的最坏情况。

在幻灯片44(或45)中给出了根据该描述的最坏情况下的未恢复访问模式示例。这并不能准确地描述最坏情况下的代码片段，但对于足够大的sliceSize来说，它是等效的。代码行是相同的。考虑到相同的读取请求(对于x，通过曲速0，在循环的第一次迭代中)，唯一的区别在于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。第三线程将读取位置2*sliceSize等。这些位置由sliceSize距离分隔。即使sliceSize仅为2，该模式的效率仍然较低，因为存储器控制器现在必须请求两倍多的行或段，以满足warp 0上的该特定读取周期。如果sliceSize足够大，内存控制器必须为每个线程请求一个唯一的行或段，这是最坏的模式。

作为最后一点，可以对"快速分析"进行有用的观察：

  • 对于大多数情况下的最佳合并访问，我们希望确保内存索引的计算不涉及threadIdx.x与1以外的任何量的乘积
  • 相反，如果我们可以证明在任何给定的索引计算中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中的相邻访问通常对共享内存也是最优的(但这不是共享内存唯一可能的最优模式)。

这里已经链接的演示将对这个主题进行更全面的处理，网络上的许多其他演示和处理也是如此。

相关文章：