c++ - 使用CUDA对两个数组求和

Question

我在学习this guide的同时正在学习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倍(22毫秒对3毫秒)。我以为，通过在连续的片上访问它们，它将执行相同或更快的操作。

我正在使用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上，所有指令都在整个warp范围内执行。因此，当warp中的一个线程正在读取内存中的位置时，warp中的所有线程都正在从内存中读取。粗略地说，最佳模式是经线中的所有线程都从相邻位置读取时。这导致以下情况:GPU内存 Controller 在检查特定时间段内warp中每个线程请求的内存地址后，可以将合并地址在一起，从而导致需要从内存请求的行数最少。高速缓存(或要从DRAM请求的最小段数)。

在幻灯片36(或37)here上以图形方式描绘了这种情况。

在您的第一个代码段中表示了100％合并的大小写。从全局内存读取的示例在这里:

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

让我们考虑循环的第一遍，并考虑第一次扭曲的情况(即线程块中的前32个线程)。在这种情况下，i由threadIdx.x给出。因此，线程0的索引为0，线程1的索引为1，依此类推。因此，每个线程都在读取全局内存中的相邻位置。假设我们错过了所有缓存，这将转化为DRAM读取请求，并且存储器 Controller 可以针对DRAM中的段(或等效地，针对缓存中的行)生成最少数量的请求(更精确地说:事务)。从“总线带宽利用率”为100％的角度来看，这是最佳的。在该读取周期中，请求的每个字节实际上都被扭曲中的线程使用。

“不公开”访问通常可以指任何不符合以上描述的情况。转化为上述更细粒度的“总线带宽利用率”数字，根据具体情况和不同情况，非强制访问的程度可能有所不同，从略低于100％的最佳情况到最坏的情况是12.5％或3.125％。 GPU。

在幻灯片44(或45)here中给出了根据此描述的最坏情况的不分时段访问模式示例。这并不能完全描述您的最坏情况的代码段，但是对于足够大的sliceSize来说，它是等效的。代码行是相同的。考虑到相同的读取请求(对于x，在循环的第一次迭代中，对于warp 0，为warp 0)，唯一的区别在于i在整个warp中采用的值:

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，该模式的效率仍然较低，因为内存 Controller 现在必须请求两倍的行数或段数才能满足在warp 0上的特定读取周期。如果sliceSize足够大，则内存 Controller 必须请求a每个线程的唯一行或段，这是最坏的情况。

作为最后的说明/要点，可以对“快速分析”做出有益的观察:

在大多数情况下可实现最佳的合并访问，我们要确保内存索引的计算不会使而不是涉及到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％合并模式也将是共享内存读/写的最佳模式，这通常是正确的。换句话说，扭曲中的相邻访问通常对于共享内存也是最佳的(但它不是共享内存的唯一可能的最佳模式)。

已经链接到here的演示文稿将对该主题进行更全面的处理，网络上的许多其他演示文稿和处理也将提供该处理。