http://blog.csdn.net/bendanban/article/details/7669624
總有些童鞋想知道怎么在CUDA中使用二維數(shù)組（[M][N]這種類型），其實這個很簡單，不過你要完全弄明白，必須對指針，地址等概念非常清楚才行。寫這篇博客解決下大家這個問題：
1、首先講述一下在一般C語言中如何使用二維數(shù)組。
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實個二維數(shù)組變量了，你可以在for循環(huán)之后arr[i][j]的方式使用它。
2、告訴你如何在CUDA中使用二維數(shù)組可以類比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中使用二維數(shù)組的幾點說明：
1）da是一個二維變量，一定更不可以在33行的時候把ha改成da！一定要記住顯存和內存是相互獨立的，主機端的程序不可以直接操作顯存！必須通過CUDA 提供的API函數(shù)來操作！
2）注意在內存申請時強制類型轉換(void**)（&），怎么把***的變量轉成**了！！這主要是API借口決定的，最好自己顯式轉換格式，避免不必要的麻煩。
3）看見數(shù)據(jù)拷貝的函數(shù)了嗎，類型、類型、還是類型。
4）別忘了釋放內存和顯存！看見沒，還是類型。
5）很希望這篇博客能幫到大家，可是我真的不推薦大家在GPU上使用二維數(shù)組！真的！！為什么呢？終歸是效率惹的禍！顯存的訪問總是慢的。二維訪存，可是連續(xù)訪問了兩次啊。要是老這樣做，不但執(zhí)行效率低，而且寫代碼也慢。如果對內存的概念不熟悉，千萬別趟這趟渾水。看懂這段代碼，就當是學習一下或者理解下內存、顯存與內存獨立的概念和規(guī)則吧。
附上執(zhí)行結果：
