背景:
项目开发过程中,有小伙伴(姑且称为小甲)写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本身二维数组申请的方式,具体做法,不知道的小伙伴,再进行查询吧...