总有些童鞋想知道怎么在CUDA中使用二维数组（[M][N]这种类型），其实这个很简单，不过你要完全弄明白，必须对指针，地址等概念非常清楚才行。写这篇博客解决下大家这个问题：
1、首先讲述一下在一般C语言中如何使用二维数组。
int r, c;
int **arr = (int**)malloc(ROWS*sizeof(int*));
int *data = (int*)malloc(COLS*ROWS*sizeof(int));
for (r = 0; r < ROWS; r++)
{
arr[r] = data + r*COLS;
}
free(arr);
free(data);
代码中的arr实个二维数组变量了，你可以在for循环之后arr[i][j]的方式使用它。
2、告诉你如何在CUDA中使用二维数组可以类比1中的方法，不过你要清楚几点，这几点在代码之后说明。
#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#define ROWS 32
#define COLS 16
#define CHECK(res) if(res!=cudaSuccess){exit(-1);}
__global__ void Kerneltest(int **da, unsigned int rows, unsigned int cols)
{
unsigned int row = blockDim.y*blockIdx.y + threadIdx.y;
unsigned int col = blockDim.x*blockIdx.x + threadIdx.x;
if (row < rows && col < cols)
{
da[row][col] = row*cols + col;
}
}
int main(int argc, char **argv)
{
int **da = NULL;
int **ha = NULL;
int *dc = NULL;
int *hc = NULL;
cudaError_t res;
int r, c;
bool is_right=true;
res = cudaMalloc((void**)(&da), ROWS*sizeof(int*));CHECK(res)
res = cudaMalloc((void**)(&dc), ROWS*COLS*sizeof(int));CHECK(res)
ha = (int**)malloc(ROWS*sizeof(int*));
hc = (int*)malloc(ROWS*COLS*sizeof(int));
for (r = 0; r < ROWS; r++)
{
ha[r] = dc + r*COLS;
}
res = cudaMemcpy((void*)(da), (void*)(ha), ROWS*sizeof(int*), cudaMemcpyHostToDevice);CHECK(res)
dim3 dimBlock(16,16);
dim3 dimGrid((COLS+dimBlock.x-1)/(dimBlock.x), (ROWS+dimBlock.y-1)/(dimBlock.y));
Kerneltest<<<dimGrid, dimBlock>>>(da, ROWS, COLS);
res = cudaMemcpy((void*)(hc), (void*)(dc), ROWS*COLS*sizeof(int), cudaMemcpyDeviceToHost);CHECK(res)
for (r = 0; r < ROWS; r++)
{
for (c = 0; c < COLS; c++)
{
printf("%4d ", hc[r*COLS+c]);
if (hc[r*COLS+c] != (r*COLS+c))
{
is_right = false;
}
}
printf("\n");
}
printf("the result is %s!\n", is_right? "right":"false");
cudaFree((void*)da);
cudaFree((void*)dc);
free(ha);
free(hc);
getchar();
return 0;
}
在CUDA中使用二维数组的几点说明：
1）da是一个二维变量，一定更不可以在33行的时候把ha改成da！一定要记住显存和内存是相互独立的，主机端的程序不可以直接操作显存！必须通过CUDA 提供的API函数来操作！
2）注意在内存申请时强制类型转换(void**)（&），怎么把***的变量转成**了！！这主要是API借口决定的，最好自己显式转换格式，避免不必要的麻烦。
3）看见数据拷贝的函数了吗，类型、类型、还是类型。
4）别忘了释放内存和显存！看见没，还是类型。
5）很希望这篇博客能帮到大家，可是我真的不推荐大家在GPU上使用二维数组！真的！！为什么呢？终归是效率惹的祸！显存的访问总是慢的。二维访存，可是连续访问了两次啊。要是老这样做，不但执行效率低，而且写代码也慢。如果对内存的概念不熟悉，千万别趟这趟浑水。看懂这段代码，就当是学习一下或者理解下内存、显存与内存独立的概念和规则吧。
附上执行结果：
