CUDA C求和2D数组的1维并返回

Question

我是GPU编程的新手(而且在C中相当生疏)所以这可能是一个相当基本的问题,我的代码中有一个明显的错误.我想要做的是采用二维数组并找到每一行的每列的总和.所以,如果我有一个包含以下内容的2D数组:

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编程指南(第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;

然而,还有更多话要说.

1. 每当您遇到CUDA代码时遇到问题,都应该使用正确的cuda错误检查.上面的第二个问题肯定会导致内存访问错误,您可能已经通过错误检查获得了一些提示.您还应该运行您的代码,cuda-memcheck这些代码将执行极其严格的边界检查,并且它肯定会捕获您的内核正在进行的越界访问.

2. 我认为你可能会对内核启动语法感到困惑:<<<ROW, COL>>> 你可能会认为这会映射到2D线程坐标(我只是猜测,因为你threadIdx.y在内核中使用它没有意义.)但是第一个参数是数字要启动的块,第二个是每个块的线程数.如果您为这两者提供标量(如您所示),您将启动一维1D线程块网格,并且您的.y变量实际上没有意义(用于索引).因此,一个要点是threadIdx.y在这种设置中没有做任何有用的事情(它始终为零).

3. 为了解决这个问题,我们可以在本回答的开头列出第一个更改.请注意,当我们启动3个块时,每个块都将具有唯一性,blockIdx.x因此我们可以将其用于索引,并且我们必须将其乘以数组的"宽度"以生成正确的索引.

4. 由于第二个参数是每个块的线程数,因此您对C的索引也没有意义.C只有3个元素(这是明智的),但每个块有10个线程,并且在每个块中线程尝试索引到C中的"前10个"位置(块中的每个线程都有唯一值threadIdx.x)但是之后在前三个位置,C中没有额外的存储空间.

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

6. 为了解决这个问题,规范的答案是使用并行缩减.这是一个复杂的话题,有关于它在这里的SO标签很多问题,所以我不会试图掩盖它,但你指出,有一个很好的教程在这里与陪同一起CUDA示例代码,你可以研究.例如,如果要查看矩阵行的平行缩减,可以查看此内容问题/答案.它恰好执行最大减少而不是减少总和,但差异很小.您也可以使用其他答案中建议的原子方法,但这通常不被视为"高性能"方法,因为原子操作的吞吐量比普通CUDA内存带宽可实现的更有限.

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