CUDA之二维数组分配内存及初始化

       在GPU中，对一维数组分配内存使用的是cudaMalloc函数，但是对于二维数组，使用cudaMalloc来分配内存并不能得到最好的性能。因为对于2D内存，对齐是一个很重要的性质，cudaMallocPitch函数能够保证分配的内存是合理对齐的，满足物理上的内存访问，因此可以保证对行访问时具有最优的效率。对数组进行初始化应当使用cudaMemset2D，进行内存赋值应当使用cudaMemcpy2D来实现。

#include<cuda_runtime.h>
#include<device_launch_parameters.h>
#include<device_functions.h>

#define M 33000
#define N 16

__gloabal__ void build_transMat(float* d_transMat, size_t pitch)
{
    int count = 1;
    for(int j = blockIdx.y * blockDim.y + threadIdx.y; j < M; j += blockDim.y * gridDim.y)
    {
        float* row_d_transMat = (float*)((char*)d_transMat + j * pitch);
        for(int i = blockIdx.x * blockDim.x + threadIdx.x; i < N; i += blockDim.x * gridDim.x)
        {
            row_d_transMat[i] = count;
            count++;
        }
    }
}

int main(int argc, char* argv[])
{
    float* d_transMat;
    float* transMat;
    size_t pitch;
    transMat = (float*)malloc(sizeof(float)*M*N)；
    cudaMallocPitch(&d_transMat, &pitch, sizeof(float) * N, M);
    cudaMemset2D(d_transMat, pitch, 0, sizeof(float) * N, M);
    build_transMat<<<blockSize, threadSize>>>(d_transMat, pitch);
    cudaMemcpy2D(transMat, sizeof(float) * N, d_transMat, pitch, sizeof(float) * N, M, cudaMemcpyDeviceToHost);
    for(int i=0; i<N; i++)
        cout << transMat[i] << endl;
    cudaFree(d_transMat);
    free(transMat);
    return 0;
}

上面的代码就是使用cudaMallocPitch来为一个M行N列的矩阵分配GPU内存空间，pitch实际上就是指一行的内存大小（sizeof(float)*N）。

       当我们使用cudaMalloc的时候，分配的是线性内存，类似于C语言中的malloc函数，连续的内存空间，从上一个元素访问到下一个相邻元素的代价比较小。如果是一个100*100的二维数组，我们依旧使用cudaMalloc来分配10000个内存空间的话，那我们访问某一行就要遍历前面所有的元素去访问，为了减小访问单行的代价，我们希望每一行的起始地址与第一行的地址是对齐的。同时，如果数组在GPU的共享内存中，通常数组会被划分到几个不同的bank中，这样有多个线程访问时就会访问到不同的bank，如果我们希望每一行可以被并行访问的话，就需要保持地址对齐。

        cudaMallocPitch所做的事情是：首先分配第一行的空间，并且检查它的总字节数是否是128的倍数，如果不是，就再多分配几个空余空间，使得总大小为128的倍数，这样一行的大小（包括补齐部分）就是一个pitch，然后以此类推分配其他行。最后，分配的总内存要大于实际所需的内存。因此，现在我们访问某一行的元素时，不是按照原来的a[i*row+j]，而是使用a[i*pitch+j]来访问。因此，使用cudaMallocPitch时，一定要返回pitch，才能访问二维数组的某个元素。

        当我们进行二维内存复制时，如果直接使用cudaMemcpy不仅复制了数组的元素，同时也复制了补齐的内存。但是，我们希望的只是复制二维数组的元素部分，就可以使用cudaMemcpy2D来使用，值复制有效元素，跳过补齐的内存。