Неправильный просмотр псевдокода CUDA?

Question

Я пытаюсь реализовать псевдокод операции prefix-sum (scan), приведенной в документации CUDA.Результаты, которые я получаю, абсолютно неверны.Я пересмотрел свой код сто раз, но у меня все еще были проблемы с этим.Вот псевдокод, данный CUDA:

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

И ядро ​​CUDA, которое я до сих пор кодировал:

// CUDA Kernel
__global__ void
prefixSumCUDA(int *a, size_t n)
{

int tId = threadIdx.x;

for (int offset = 1; offset < n; offset *= 2) {
    if (tId >= pow((float)2, offset)) {
        int temp = tId - pow((float)2, offset - 1);
        a[tId] += a[temp];
    }
}
}

Пожалуйста, дайте мне знать, если я допускаю какие-либо ошибки здесь.Я знаю, что эта реализация в значительной степени зависит от размера блоков и сеток.Таким образом, я предоставлю свой вызов ядра здесь:

// Kernel launch
prefixSumCUDA << <1, 32 >> > (d_A, n);

Входной массив имеет целочисленный тип из 8 элементов:

[-] array: 1, 2, 3, 4, 5, 6, 7, 8

И результат ядра CUDA будет следующим:

[-] array: 1, 2, 5, 7, 14, 18, 22, 26

Спасибо за любую помощь заранее!

Answer 1

Я решил проблему, реализовав это другим способом.Смещение лучше начинать с 0, а не с 1. В результате получается следующий код.

__global__ void
prefixSumCUDA(int *a, size_t n)
{

int tId = threadIdx.x;

int end = ceil(log2((float)n));

for (int offset = 0; offset < end; offset++) {
    if (tId >= (1 << offset)) {
        a[tId] += a[tId - (1 << offset)];
    }
}
}