二维数组和二维指针在CUDA中的应用

 CUDA是一个线程网络，我特别想弄清楚的一件事情是，主机如果是个二维的数组，传到设备中，是否还是可以用二维数组来表示呢？很多例子都是将二维的数组映射
 到一个一维指针变量中去。但是我就是还想在设备中也用arr[][]的形式来找到我想要的那个元素，可以吗？肯定是可以的。
 方案一：栈
 定义arr[2][10]，直接用地址传进去。但是栈的空间很小，在CPU中能声明的数组就不大。所以，这里不讨论这个方案。
 方案二：堆上new出个二维指针
 结合CPU来说：
 首先，在CPU中，二维指针和二维数组的一个映射关系可以用一个拷贝实现，代码入下：

……
 21     // b[2][10]
 22     int **b= (int**)malloc(2*sizeof(int*));
 23     int *data = (int*)malloc(2*10*sizeof(int));
 24     for(int i=0;i<20;i++)
 25         data[i] = i;
 26     // 将数据复制到二维的指针变量中，是数据赋值吗？还是地址哦？我的观点会在最后给出。
 27     for(int i=0;i<2;i++)
 28         b[i] = 10*i + data;

 之后，二维指针变量b，就可以用b[][]的二维数组访问形式来访问你想要的那个元素了。 
 30     for(int i=0;i<2;i++){
 31         for(int j=0;j<10;j++){
 32             cout << b[i][j] << " ";
 33         }
 34         cout << endl;
 35     }
……

接着，我想看看CPU和GPU上有什么不同。 在主机函数申请号b和data的空间后，22、23行代码，我们开辟对应的设备内存 int **dev_b; int *dev_data; cudaMalloc((void**)(&dev_b), ROWS*sizeof(int*)); cudaMalloc((void**)(&dev_data), ROWS*COLS*sizeof(int)); 我们需要用27、28行代码来赋值。

如何拷贝主机和设备呢？ cudaMemcpy((void*)(dev_b), (void*)(b), 2*sizeof(int*), cudaMemcpyHostToDevice); 要注意，这里将二维指针的内容进行拷贝，因为，我们的目的就是想在核函数中用arr[][]的形式访问元素。

下面是核函数了：

__global__ void kernelfun(int **dev_b)//我们就作个自增1的事情吧
{
    unsigned int row = blockDim.y*blockIdx.y + threadIdx.y;  
    unsigned int col = blockDim.x*blockIdx.x + threadIdx.x;  
    if (row < 2 && col < 10)  
    {  
        dev_b[row][col] += 1;  
    }
}

解释一下，主机函数我们用了二维的线程网络：

dim3 dimBlock(16,16); 

dim3 dimGrid((COLS+dimBlock.x-1)/(dimBlock.x), (ROWS+dimBlock.y-1)/(dimBlock.y)); 

Kerneltest<<<dimGrid, dimBlock>>>(dev_b); 

悲哀的是，我们这里有2*10个元素，只用了一个block中的2*10个线程……

最后，我们现在要将计算结果拷贝出来，就算完成要做的事情了，是不是拷贝dev_b和dev_data都可以呢？留给自己一个作业题，反正初步一看，这两个东西是同一个， 因为赋值的是地址。对甲操作实质上就是对乙操作。所以，其实数据上来说是一回事。