使用CUDA的前缀和

我很难理解简单前缀和的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”变量?他们是做什么的?因为这里最多只使用一个块和1024个线程,所以我们只能使用threadId。x指定块中的元素?
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,如果有人能回答我的问题,我将不胜感激。。。