CUDA C对二维数组的一维求和并返回

Question

我是 GPU 编程的新手（在 C 语言中相当生疏），所以这可能是一个相当基本的问题，我的代码中有一个明显的错误。我想要做的是获取一个二维数组并找到每一行的每一列的总和。因此，如果我有一个包含以下内容的二维数组：

0 1 2 3 4 5 6 7 8 9
0 1 2 3 4 5 6 7 8 9
0 2 4 6 8 10 12 14 16 18

我想得到一个包含以下内容的数组：

45
45
90

到目前为止，我的代码没有返回正确的输出，我不知道为什么。我猜这是因为我没有正确处理内核中的索引。但可能是我没有正确使用内存，因为我从一个过度简化的一维示例中调整了它，CUDA Programming Guide（第 3.2.2 节）为初学者在 1和二维数组。

我的错误尝试：

#include <stdio.h>
#include <stdlib.h>


// start with a small array to test
#define ROW 3
#define COL 10

__global__ void collapse( int *a, int *c){
    /*
       Sum along the columns for each row of the 2D array.
    */
    int total = 0;
    // Loop to get total, seems wrong for GPUs but I dont know a better way
    for (int i=0; i < COL; i++){
        total = total + a[threadIdx.y + i];
    }
    c[threadIdx.x] = total;

}

int main( void ){
    int array[ROW][COL];      // host copies of a, c
    int c[ROW];
    int *dev_a;      // device copies of a, c (just pointers)
    int *dev_c;

    // get the size of the arrays I will need
    int size_2d = ROW * COL * sizeof(int);
    int size_c = ROW * sizeof(int);

    // Allocate the memory
    cudaMalloc( (void**)&dev_a, size_2d);
    cudaMalloc( (void**)&dev_c, size_c);

    // Populate the 2D array on host with something small and known as a test
    for (int i=0; i < ROW; i++){
        if (i == ROW - 1){
            for (int j=0; j < COL; j++){
                array[i][j] = (j*2);
                printf("%i ", array[i][j]);
            }
        } else {
            for (int j=0; j < COL; j++){
                array[i][j] = j;
                printf("%i ", array[i][j]);
            }
        }
        printf("\n");
    }

    // Copy the memory
    cudaMemcpy( dev_a, array, size_2d, cudaMemcpyHostToDevice );
    cudaMemcpy( dev_c, c, size_c, cudaMemcpyHostToDevice );

    // Run the kernal function
    collapse<<< ROW, COL >>>(dev_a, dev_c);

    // copy the output back to the host
    cudaMemcpy( c, dev_c, size_c, cudaMemcpyDeviceToHost );

    // Print the output
    printf("\n");
    for (int i = 0; i < ROW; i++){
        printf("%i\n", c[i]);
    }

    // Releasae the memory
    cudaFree( dev_a );
    cudaFree( dev_c );
}

输出：

0 1 2 3 4 5 6 7 8 9
0 1 2 3 4 5 6 7 8 9
0 2 4 6 8 10 12 14 16 18

45
45
45

Answer 1

你是对的，这是一个索引问题。如果你替换这个，你的内核会生成一个正确的答案：

    total = total + a[threadIdx.y + i];

用这个：

    total = total + a[blockIdx.x*COL + i];

还有这个：

c[threadIdx.x] = total;

用这个：

c[blockIdx.x] = total;

不过还有很多话要说。

1. 任何时候您在使用 CUDA 代码时遇到问题，都应该使用 proper cuda error checking。上面的第二个问题肯定会导致内存访问错误，您可能已经通过错误检查得到了提示。您还应该使用cuda-memcheck 运行您的代码，这将执行非常严格的边界检查工作，并且它肯定会捕获您的内核正在进行的越界访问。

2. 我认为您可能对内核启动语法感到困惑：<<<ROW, COL>>> 您可能认为这会映射到 2D 线程坐标（我只是在猜测，因为您在具有它的内核中使用了 threadIdx.y没有意义。）但是第一个参数是要启动的blocks的数量，第二个是每个block的线程数。如果您为这两个提供标量（正如您所拥有的），您将启动一维线程块的一维网格，并且您的 .y 变量将没有真正意义（对于索引）。所以一个要点是threadIdx.y 在这个设置中没有做任何有用的事情（它总是为零）。

3. 要解决此问题，我们可以进行此答案开头列出的第一个更改。请注意，当我们启动 3 个块时，每个块将有一个唯一的blockIdx.x，因此我们可以将其用于索引，并且我们必须将其乘以数组的“宽度”以生成正确的索引。

   李>

4. 由于第二个参数是每个块的线程数，因此您对 C 的索引也没有意义。 C 只有 3 个元素（这是明智的），但每个块有 10 个线程，并且在每个块中，线程试图索引到 C 中的“前 10 个”位置（块中的每个线程都有一个唯一值 threadIdx.x ) 但是在前 3 个位置之后，C 中就没有多余的存储空间了。

5. 现在可能是最大的问题。 块中的每个线程都在循环中做完全相同的事情。您的代码不会区分线程的行为。您可以通过这种方式编写给出正确答案的代码，但从性能的角度来看这是不明智的。

6. 要解决最后一个问题，规范的答案是使用并行归约。这是一个涉及的话题，在 SO 标签上有很多关于它的问题，所以我不会尝试覆盖它，但向您指出，有一个很好的教程 here 以及随附的 CUDA sample code你可以学习。例如，如果您想查看矩阵行的并行缩减，可以查看此question/answer。它恰好是执行最大减少而不是总和减少，但差异很小。您也可以使用另一个答案中建议的原子方法，但这通常不被视为“高性能”方法，因为原子操作的吞吐量比普通 CUDA 内存带宽所能达到的更有限。

您似乎也对 CUDA 内核执行模型感到困惑，因此继续阅读编程指南（您已经链接）是一个很好的起点。