使用 CUDA 的前缀和

Question

我无法理解幼稚前缀总和的 cuda 代码。

这是代码来自https://developer.nvidia.com/gpugems/GPUGems3/gpugems3_ch39.html 在示例 39-1（朴素扫描）中，我们有这样的代码：

 __global__ void scan(float *g_odata, float *g_idata, int n)
    {
    extern __shared__ float temp[]; // allocated on invocation
    int thid = threadIdx.x;
    int pout = 0, pin = 1;
    // Load input into shared memory.
    // This is exclusive scan, so shift right by one
    // and set first element to 0
    temp[pout*n + thid] = (thid > 0) ? g_idata[thid-1] : 0;
 __syncthreads();

 for (int offset = 1; offset < n; offset *= 2)
  {
    pout = 1 - pout; // swap double buffer indices
    pin = 1 - pout;
    if (thid >= offset)
      temp[pout*n+thid] += temp[pin*n+thid - offset];
    else
      temp[pout*n+thid] = temp[pin*n+thid];
    __syncthreads();
  }
  g_odata[thid] = temp[pout*n+thid1]; // write output
}

我的问题是

1. 为什么我们需要创建共享内存临时？
2. 为什么我们需要“pout”和“pin”变量？他们在做什么？既然我们这里只用了一个block，最多1024个线程，那我们能不能只用threadId.x来指定block中的元素呢？
3. 在CUDA中，我们使用一个线程做一个添加操作吗？如果我使用 for 循环（给定一个线程用于数组中的一个元素，循环 OpenMP 中的线程或处理器），一个线程是否可以在一次迭代中完成？
4. 我之前的两个问题可能看起来很幼稚......我认为关键是我不明白上述实现与伪代码之间的关系如下：

for d = 1 to log2 n do for all k in parallel do if k >= 2^d then x[k] = x[k – 2^(d-1)] + x[k]

这是我第一次使用 CUDA，如果有人能回答我的问题，我将不胜感激......

Answer 1

1- 将内容放入共享内存 (SM) 并在那里进行计算比使用全局内存更快。加载 SM 后同步线程很重要，因此需要 __syncthreads。

---

2- 这些变量可能是为了澄清算法中的顺序颠倒。它只是用于切换某些部分：

temp[pout*n+thid] += temp[pin*n+thid - offset];

第一次迭代； pout = 1 和 pin = 0。第二次迭代； pout = 0 和 pin = 1。 它在奇数迭代中偏移 N 数量的输出，并在偶数迭代中偏移输入。回到您的问题，您无法使用 threadId.x 实现相同的目标，因为它不会在循环内改变。

---

3 & 4 - CUDA 执行线程来运行内核。这意味着每个线程单独运行该代码。如果您查看伪代码并与 CUDA 代码进行比较，您已经将外部循环与 CUDA 并行化了。所以每个线程都会在内核中运行循环，直到循环结束，并在写入全局内存之前等待每个线程完成。

---

希望对你有帮助。