cuda二维数组内存分配和数据拷贝

2016-04-20 10:54 138人阅读 评论(0) 收藏 举报 分类： 机器学习（11） 人工智能（9）

版权声明：本文为博主原创文章，允许转载。

我们看一个例子，如何对矩阵进行分配显卡内存以及元素赋值操作。通常来讲，在GPU中分配内存使用的是cudaMalloc函数，但是对于二维或者三维矩阵而言，使用cudaMalloc来分配内存并不能得到最好的性能，原因是对于2D或者3D内存，对齐是一个很重要的性质，而cudaMallocPitch或者cudaMalloc3D这两个函数能够保证分配的内存是合理对齐的，满足物理上的内存访问，因此可以确保对行访问时具有最优的效率，除此之外，对于数组内存的复制应当使用cudaMemcpy2D和cudaMemcpy3D来实现。而对于一维数组，使用cudaMalloc以及cudaMemcpy即可满足使用需求。下面我们通过一段代码来认识一下cudaMallocPitch这个函数。

?

1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35

#define N 11 #define M 3 #define GridSize 16 #define BlockSize 16 #include<iostream> usingnamespacestd;   __global__voidkernel(float* d_matrix, size_tpitch) {     intcount = 1;     for(intj = blockIdx.y * blockDim.y + threadIdx.y; j < N; j += blockDim.y * gridDim.y)     {         float* row_d_matrix = (float*)((char*)d_matrix + j*pitch);         for(inti = blockIdx.x * blockDim.x + threadIdx.x; i < M; i += blockDim.x * gridDim.x)         {             row_d_matrix[i] = count;             count++;         }     } }   intmain() {     float*d_matrix;     float*dc_matrix;     dc_matrix = (float*)malloc(sizeof(float)*M*N);     size_tpitch;     cudaMallocPitch(&d_matrix,&pitch,M*sizeof(float),N);     kernel<<<GridSize,BlockSize>>>(d_matrix,pitch);     cudaMemcpy2D(dc_matrix, M * sizeof(float), d_matrix, pitch, M * sizeof(float), N, cudaMemcpyDeviceToHost);     for(inti=0;i<M*N;i++)         cout<<dc_matrix[i]<<endl;     cudaFree(d_matrix);     free(dc_matrix);     return0; }

上面这段代码就是使用cudaMallocPitch来为一个N行M列的矩阵分配GPU内存空间，这里的pitch实际上就是指一行的内存大小，在这里就是M*sizeof(float)，当然有人可能会问pitch的作用是什么，一般来说，当我们使用cudaMalloc的时候，它分配的是线性内存，类似于C语言中的malloc函数，连续的内存空间，从上一个元素到访问下一个相邻元素的代价比较小。但是如果我们希望得到一个100*100的二维数组的话，如果我们依旧使用cudaMalloc来分配10000个内存空间的话，那我们访问某一行的话我们就要遍历前面的所有元素去访问，为了减小访问单行的代价，我们希望我们的每一行起始地址与第一行的地址是对齐的。同时，如果数组是在GPU的共享内存中，通常它会被划分到几个不同的bank中，这样有多个线程访问时就会访问到不同的bank，如果我们希望我们的每一行可以被并行访问的话，我门就需要保持地址对齐。cudaMallocPitch所做的事情就是：首先分配第一行的空间，并且检查它的总字节数是否是128的倍数，如果不是的话，就再多分配几个空余空间，使得总大小为128的倍数，这个一行的大小（包括补齐部分）就是一个pitch，然后以此类推分配其他行。最后，分配的总内存要大于实际所需的内存。因此，现在我们访问某一行的某个元素时，就不是按照原来的想法a[row*i+j]而是使用a[pitch*i+j]来访问，就是这个原理。因此，我们使用cudaMallocPitch的时候，一定要返回pitch的，只有这样，我们才能访问二维数组的某个元素。

同理，当我们需要进行二维内存复制时，假如我们希望将这个二维数组从device复制到host的时候，如果直接使用cudaMemcpy则不仅复制了数组的元素，同时也复制了补齐的内存，这是我们不希望得到的。即我们想告诉GPU我们只希望复制二维数组的元素部分，这个时候就可以使用cudaMemcpy2D来实现这个功能了，只复制有效元素，跳过补齐的内存。