不明白为什么CUDA中列添加比行添加快

Question

我从 CUDA 开始，编写了两个内核进行实验。Whey 都接受 3 个指向 n*n（矩阵仿真）和 n 数组的指针。

__global__
void th_single_row_add(float* a, float* b, float* c, int n) {
  int idx = blockDim.x * blockIdx.x * n + threadIdx.x * n;
  for (int i = 0; i < n; i ++) {
    if (idx + i >= n*n) return;
    c[idx + i] = a[idx + i] + b[idx + i];
  }
}

__global__
void th_single_col_add(float* a, float* b, float* c, int n) {
  int idx = blockDim.x * blockIdx.x + threadIdx.x;
  for (int i = 0; i < n; i ++) {
    int idx2 = idx + i * n;
    if (idx2 >= n*n) return;
    c[idx2] = a[idx2] + b[idx2];
  }
}

在th_single_row_add每个线程中对元素求和行n，在th_single_col_add每个线程求和列。n = 1000这是（1 000 000 个元素）的简介

986.29us  th_single_row_add(float*, float*, float*, int)
372.96us  th_single_col_add(float*, float*, float*, int)

如您所见，列求和速度快了三倍。我认为因为在column变体中循环中的所有索引彼此相距很远，所以它应该更慢，我错在哪里？

Answer 1

CUDA 中的线程不会单独运行，它们以32 个线程的扭曲形式分组在一起。这 32 个线程（通常）同步执行。向一个线程发出的指令会在同一时钟周期内同时向所有 32 个线程发出。

如果该指令是读取存储器的指令（例如），则可能需要/请求最多32次独立读取。满足这些读取操作所需的确切地址模式由您编写的代码决定。如果这些地址在内存中全部“相邻”，那么这将是一次有效的读取。如果这些地址以某种方式“分散”在内存中，那么读取的效率将很低，并且速度会更慢。

刚刚描述的这个基本概念在 CUDA 中称为“合并”访问。您的列求和情况允许跨扭曲进行合并访问，因为扭曲中每个线程生成的地址位于相邻列中，并且位置在内存中相邻。您的行求和案例打破了这一点。经线中每个线程生成的地址不相邻（它们是“列”的，按数组的宽度彼此分隔），因此不是“合并”的。

性能差异是由于内存访问效率的差异造成的。

您可以通过研究 CUDA 优化的介绍性处理来研究有关 CUDA 中合并行为的更多信息，例如此处特别是幻灯片 44-54。