Cuda手撕二维数组

背景：

    项目开发过程中，有小伙伴（姑且称为小甲）写cuda程序一直运行出错，且百思不得其解，反复查看逻辑，感觉没有问题，一起review代码发现，其cuda逻辑中使用了二维指针，且使用错误，cuda二维指针的使用，初学者如果不使用cuda本身二维数组的申请，直接使用cudaMalloc，及其容易出错，因此写该篇文章，以供参考。

1. 一维数组的申请

我们从一维数据说起，大家都知道，一维数组申请及拷贝形式：

（代码手敲，有可能有语法错误，请见谅）

const int i_size = 5;

int *dev_a = 0;
cudaMalloc((void**)&dev_a, i_size * sizeof(int));

int host_a[i_size] = {1,2,3,4,5};
cudaMemcpy(dev_a, host_a, i_size * sizeof(int), cudaMemcpyHostToDevice);

其中dev_a最开始定义的时候是主机地址，但是我们最终是要申请一个Gpu地址，怎么操作呢？即上面代码中的cudaMalloc，我们发现，执行cudaMalloc的时候，第一个参数是使用的dev_a的地址作为参数，传入cudaMalloc的，为什么?

原因是，cudaMalloc中，会通过寻址，将dev_a的值，修改为显存的首地址。

类似上图了，传给cudaMalloc的是Ptr的值，然后cudaMalloc通过*Ptr，将dev_a的值，从指向null，修改为指向一片显存。

OK，一维数据讲完了，如果有不清楚的，请留言沟通，下面我们引入前面说到的项目中的二维数组问题

2. 二维数组的申请

 先来看下小甲是怎么写的吧（代码手敲，有可能有语法错误，请见谅）

const int i_size = 100;
int *dev_a[4] = {0， 0，0，0};

for (int i = 0; i < 4; ++i)
{
    cudaMalloc((void**)&dev_a[i], i_size * sizeof(int));
}

...

ker<<<1, 1>>>(dev_a);

写到这里，大家有发现问题吗？

假装思考ing...

OK，继续

    如果我们希望dev_a能够作为正常的一个指针传给核函数进行逻辑运行，那应该怎么做，从前面第一节一维数组可知，应该是要将dev_a的值，修改为指向显存的地址。那上面的代码做的事情是什么？

上面代码只是将dev_a指向的4个null修改为了指向一片显存的首地址，但是dev_a并没有进行处理，仍然是主机端地址值：

 

 因此，直接将dev_a传到核函数中进行访问是错误的。。。那该怎么办，下面是一个直接使用cudaMalloc的实例，可直接运行：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size);

__global__ void addKernel(int *c, int **a, int **b, int size)
{
    for (int k = 0; k < size; ++k)
    {
        for (int kk = 0; kk < 4; ++kk)
        {
            c[k] += a[kk][k] + b[kk][k];
        }
    }
}

int main()
{
    const int arraySize = 5;
    const int a[arraySize] = { 1, 2, 3, 4, 5 };
    const int b[arraySize] = { 10, 20, 30, 40, 50 };
    int c[arraySize] = { 0 };

    addWithCuda(c, a, b, arraySize);

    printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",
        c[0], c[1], c[2], c[3], c[4]);

    return 0;
}

// Helper function for using CUDA to add vectors in parallel.
cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size)
{
    int **dev_a = 0;
    int **dev_b = 0;
    int *dev_c = 0;

    cudaMalloc((void**)&dev_c, size * sizeof(int));
    cudaMalloc((void**)&dev_a, 4 * sizeof(int*));
    cudaMalloc((void**)&dev_b, 4 * sizeof(int*));

    int* h_a[4] = {0};
    int* h_b[4] = { 0 };

    for (int i = 0; i < 4; ++i)
    {
        cudaMalloc(&h_a[i], size * sizeof(int));
        cudaMemcpy(h_a[i], a, sizeof(int) * size, cudaMemcpyHostToDevice);

        cudaMalloc(&h_b[i], size * sizeof(int));
        cudaMemcpy(h_b[i], b, sizeof(int) * size, cudaMemcpyHostToDevice);
    }

    cudaMemcpy(dev_a, h_a, sizeof(int*) * 4, cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b, h_b, sizeof(int*) * 4, cudaMemcpyHostToDevice);

    // Launch a kernel on the GPU with one thread for each element.
    addKernel<<<1, 1>>>(dev_c, dev_a, dev_b, size);
    cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
    return 0;
}

写到这里，应该差不多了吧，当然不推荐这样处理，还是使用cuda本身二维数组申请的方式，具体做法，不知道的小伙伴，再进行查询吧...