二維數組和二維指針在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都可以呢？留給自己一個作業題，反正初步一看，這兩個東西是同一個， 因為賦值的是地址。對甲操作實質上就是對乙操作。所以，其實數據上來說是一回事。