CUDA用於並行計算非常方便，但是GPU與CPU之間的交互，比如傳遞參數等相對麻煩一些。在寫CUDA核函數的時候形參往往會有很多個，動輒達到10-20個，如果能夠在CPU中提前把數據組織好，比如使用二維數組，這樣能夠省去很多參數，在核函數中可以使用二維數組那樣去取數據簡化代碼結構。當然使用二維數據會增加GPU內存的訪問次數，不可避免會影響效率，這個不是今天討論的重點了。
舉兩個代碼栗子來說明二維數組在CUDA中的使用（親測可用）：
1. 普通二維數組示例：
輸入：二維數組A（8行4列）
輸出：二維數組C（8行4列）
函數功能：將數組A中的每一個元素加上10，並保存到C中對應位置。
這個是一個簡單的示例，以一級指針和二級指針開訪問二維數組中的數據，主要步驟如下：
（1）為二級指針A、C和一級指針dataA、dataC分配CPU內存。二級指針指向的內存中保存的是一級指針的地址。一級指針指向的內存中保存的是輸入、輸出數據。
（2）在設備端（GPU）上同樣建立二級指針d_A、d_C和一級指針d_dataA、d_dataC，並分配GPU內存，原理同上，不過指向的內存都是GPU中的內存。
（3）通過主機端一級指針dataA將輸入數據保存到CPU中的二維數組中。
（4）關鍵一步：將設備端一級指針的地址，保存到主機端二級指針指向的CPU內存中。
（5）關鍵一步：使用cudaMemcpy（）函數，將主機端二級指針中的數據（設備端一級指針的地址）拷貝到設備端二級指針指向的GPU內存中。這樣在設備端就可以使用二級指針來訪問一級指針的地址，然后利用一級指針訪問輸入數據。也就是A[][]、C[][]的用法。
（6）使用cudaMemcpy（）函數將主機端一級指針指向的CPU內存空間中的輸入數據，拷貝到設備端一級指針指向的GPU內存中，這樣輸入數據就算上傳到設備端了。
（7）在核函數addKernel（）中就可以使用二維數組的方法進行數據的讀取、運算和寫入。
（8）最后將設備端一級指針指向的GPU內存中的輸出數據拷貝到主機端一級指針指向的CPU內存中，打印顯示即可。
#include <cuda_runtime.h> #include <device_launch_parameters.h> #include <opencv2\opencv.hpp> #include <iostream> #include <string> using namespace cv; using namespace std; #define Row 8 #define Col 4 __global__ void addKernel(int **C, int **A) { int idx = threadIdx.x + blockDim.x * blockIdx.x; int idy = threadIdx.y + blockDim.y * blockIdx.y; if (idx < Col && idy < Row) { C[idy][idx] = A[idy][idx] + 10; } } int main() { int **A = (int **)malloc(sizeof(int*) * Row); int **C = (int **)malloc(sizeof(int*) * Row); int *dataA = (int *)malloc(sizeof(int) * Row * Col); int *dataC = (int *)malloc(sizeof(int) * Row * Col); int **d_A; int **d_C; int *d_dataA; int *d_dataC; //malloc device memory cudaMalloc((void**)&d_A, sizeof(int **) * Row); cudaMalloc((void**)&d_C, sizeof(int **) * Row); cudaMalloc((void**)&d_dataA, sizeof(int) *Row*Col); cudaMalloc((void**)&d_dataC, sizeof(int) *Row*Col); //set value for (int i = 0; i < Row*Col; i++) { dataA[i] = i+1; } //將主機指針A指向設備數據位置，目的是讓設備二級指針能夠指向設備數據一級指針 //A 和 dataA 都傳到了設備上，但是二者還沒有建立對應關系 for (int i = 0; i < Row; i++) { A[i] = d_dataA + Col * i; C[i] = d_dataC + Col * i; } cudaMemcpy(d_A, A, sizeof(int*) * Row, cudaMemcpyHostToDevice); cudaMemcpy(d_C, C, sizeof(int*) * Row, cudaMemcpyHostToDevice); cudaMemcpy(d_dataA, dataA, sizeof(int) * Row * Col, cudaMemcpyHostToDevice); dim3 block(4, 4); dim3 grid( (Col + block.x - 1)/ block.x, (Row + block.y - 1) / block.y ); addKernel << <grid, block >> > (d_C, d_A); //拷貝計算數據-一級數據指針 cudaMemcpy(dataC, d_dataC, sizeof(int) * Row * Col, cudaMemcpyDeviceToHost); for (int i = 0; i < Row*Col; i++) { if (i%Col == 0) { printf("\n"); } printf("%5d", dataC[i]); } printf("\n"); }
2.OpenCV中Mat數組示例
輸入：圖像Lena.jpg
輸出：圖像moon.jpg
函數功能：求兩幅圖像加權和
原理和上面一樣，流程上的差別就是輸入的二維數據是下面兩幅圖像數據，然后在CUDA中進行加權求和。
效果如下：
代碼在此，以供參考
#include <cuda_runtime.h> #include <device_launch_parameters.h> #include <opencv2\opencv.hpp> #include <iostream> #include <string> using namespace cv; using namespace std; __global__ void addKernel(uchar **pSrcImg, uchar* pDstImg, int imgW, int imgH) { int tidx = threadIdx.x + blockDim.x * blockIdx.x; int tidy = threadIdx.y + blockDim.y * blockIdx.y; if (tidx<imgW && tidy<imgH) { int idx=tidy*imgW+tidx; uchar lenaValue=pSrcImg[0][idx]; uchar moonValue=pSrcImg[1][idx]; pDstImg[idx]= uchar(0.5*lenaValue+0.5*moonValue); } } int main() { //OpenCV讀取兩幅圖像 Mat img[2]; img[0]=imread("data/lena.jpg", 0); img[1]=imread("data/moon.jpg", 0); int imgH=img[0].rows; int imgW=img[0].cols; //輸出圖像 Mat dstImg=Mat::zeros(imgH, imgW, CV_8UC1); //主機指針 uchar **pImg=(uchar**)malloc(sizeof(uchar*)*2); //輸入 二級指針 //設備指針 uchar **pDevice;//輸入 二級指針 uchar *pDeviceData;//輸入 一級指針 uchar *pDstImgData;//輸出圖像對應設備指針 //分配GPU內存 cudaError err; //目標輸出圖像分配GPU內存 err=cudaMalloc(&pDstImgData, imgW*imgH*sizeof(uchar)); //設備二級指針分配GPU內存 err=cudaMalloc(&pDevice, sizeof(uchar*)*2); //設備一級指針分配GPU內存 err=cudaMalloc(&pDeviceData, sizeof(uchar)*imgH*imgW*2); //關鍵：主機二級指針指向設備一級指針位置，這樣才能使設備的二級指針指向設備的一級指針位置 for (int i=0; i<2; i++) { pImg[i]=pDeviceData+i*imgW*imgH; } //拷貝數據到GPU //拷貝主機二級指針中的元素到設備二級指針指向的GPU位置 （這個二級指針中的元素是設備中一級指針的地址） err=cudaMemcpy(pDevice, pImg, sizeof(uchar*)*2, cudaMemcpyHostToDevice); //拷貝圖像數據(主機一級指針指向主機內存) 到 設備一級指針指向的GPU內存中 err=cudaMemcpy(pDeviceData, img[0].data, sizeof(uchar)*imgH*imgW, cudaMemcpyHostToDevice); err=cudaMemcpy(pDeviceData+imgH*imgW, img[1].data, sizeof(uchar)*imgH*imgW, cudaMemcpyHostToDevice); //核函數實現lena圖和moon圖的簡單加權和 dim3 block(8, 8); dim3 grid( (imgW+block.x-1)/block.x, (imgH+block.y-1)/block.y); addKernel<<<grid, block>>>(pDevice, pDstImgData, imgW, imgH); cudaThreadSynchronize(); //拷貝輸出圖像數據至主機，並寫入到本地 err=cudaMemcpy(dstImg.data, pDstImgData, imgW*imgH*sizeof(uchar), cudaMemcpyDeviceToHost); imwrite("data/synThsis.jpg", dstImg); }