在CUDA怎么使用二位数组（*[M][N]）_CUDA

在CUDA如何使用二位数组（**[M][N]）

总有些童鞋想知道怎么在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上使用二维数组！真的！！为什么呢？终归是效率惹的祸！显存的访问总是慢的。二维访存，可是连续访问了两次啊。要是老这样做，不但执行效率低，而且写代码也慢。如果对内存的概念不熟悉，千万别趟这趟浑水。看懂这段代码，就当是学习一下或者理解下内存、显存与内存独立的概念和规则吧。

附上执行结果：