Сумма префикса с использованием 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"? Что они делают? Поскольку здесь мы используем только один блок и максимум 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, поэтому буду признателен, если кто-нибудь ответит на мои вопросы...

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. Таким образом, каждый поток будет выполнять цикл в ядре до конца цикла и будет ждать завершения каждого потока перед записью в глобальную память.

---

Надеюсь, поможет.