CUDA C 求和二维数组的 1 维和 return

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

我目前的代码没有返回正确的输出，我不确定为什么。我猜这是因为我没有正确处理内核中的索引。但这可能是我没有正确使用内存，因为我从一个过度简化的 1 维示例中改编了它，并且 CUDA Programming Guide（第 3.2.2 节）为初学者做了一个相当大且描述不太好的跳跃介于 1 和 2 维数组之间。

我的错误尝试：

#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;

然而还有很多话要说。

任何时候您在使用 CUDA 代码时遇到问题，您都应该使用 proper cuda error checking。上面的第二个问题肯定会导致内存访问错误，您可能已经通过错误检查得到了提示。您还应该运行您的代码 cuda-memcheck ，这将执行非常严格的边界检查工作，并且它肯定会捕获您的内核正在进行的越界访问。
我认为您可能对内核启动语法感到困惑：<<<ROW, COL>>>您可能认为这映射到 2D 线程坐标（我只是猜测，因为您使用了 threadIdx.y 在内核中没有意义。）但是第一个参数是要启动的块数，第二个参数是每个块的 线程数。如果您为这两者提供标量（正如您所拥有的那样），您将启动一维线程块的一维网格，并且您的 .y 变量实际上没有意义（对于索引）。所以一个要点是 threadIdx.y 在此设置中没有做任何有用的事情（它始终为零）。
要解决此问题，我们可以进行本答案开头列出的第一项更改。请注意，当我们启动 3 个块时，每个块都有一个唯一的 blockIdx.x，因此我们可以将其用于索引，我们必须将其乘以数组的 "width" 以生成正确的索引。
由于第二个参数是每个块的线程数，因此您对 C 的索引也没有意义。 C 只有 3 个元素（这是合理的），但每个块有 10 个线程，并且在每个块中，线程试图索引到 C 中的 "first 10" 位置（块中的每个线程都有一个唯一的值 threadIdx.x) 但是在前3个位置之后，C中就没有多余的存储了。
现在可能是最大的问题。 块中的每个线程在循环中做完全相同的事情。您的代码不区分线程的行为。您可以通过这种方式编写给出正确答案的代码，但从性能的角度来看，这是不明智的。
要解决最后一个问题，规范的答案是使用并行缩减。这是一个涉及的话题，在 SO 标签上有很多关于它的问题，所以我不会试图涵盖它，但要指出有一个很好的教程 here along with the accompanying CUDA sample code that you can study. If you want to see a parallel reduction on the matrix rows, for example, you could look at this 。它恰好执行的是最大约简而不是和约简，但差异很小。您也可以使用其他答案中建议的原子方法，但这通常不被视为 "high-performance" 方法，因为原子操作的吞吐量比普通 CUDA 内存带宽可实现的吞吐量更有限。

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

CUDA C 求和二维数组的 1 维和 return

CUDA C sum 1 dimension of 2D array and return

c

arrays

indexing

cuda

nvidia