uda二維數組內存分配和數據拷貝

2016-04-20 10:54 138人閱讀 評論(0) 收藏 舉報 分類： 機器學習（11） 人工智能（9）

版權聲明：本文為博主原創文章，允許轉載。

因為cuda具有高效利用GPU進行科學計算的優勢，而人工智能的重點之一就是復雜的計算任務，因此學好GPU計算是學習AI的重點任務。這里，我們即將進行利用共享內存的矩陣運算。

我們看一個例子，如何對矩陣進行分配顯卡內存以及元素賦值操作。通常來講，在GPU中分配內存使用的是cudaMalloc函數，但是對于二維或者三維矩陣而言，使用cudaMalloc來分配內存并不能得到最好的性能，原因是對于2D或者3D內存，對齊是一個很重要的性質，而cudaMallocPitch或者cudaMalloc3D這兩個函數能夠保證分配的內存是合理對齊的，滿足物理上的內存訪問，因此可以確保對行訪問時具有最優的效率，除此之外，對于數組內存的復制應當使用cudaMemcpy2D和cudaMemcpy3D來實現。而對于一維數組，使用cudaMalloc以及cudaMemcpy即可滿足使用需求。下面我們通過一段代碼來認識一下cudaMallocPitch這個函數。

?

1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35

#define N 11 #define M 3 #define GridSize 16 #define BlockSize 16 #include<iostream> usingnamespacestd; __global__voidkernel(float* d_matrix, size_tpitch) { ????intcount = 1; ????for(intj = blockIdx.y * blockDim.y + threadIdx.y; j < N; j += blockDim.y * gridDim.y) ????{ ????????float* row_d_matrix = (float*)((char*)d_matrix + j*pitch); ????????for(inti = blockIdx.x * blockDim.x + threadIdx.x; i < M; i += blockDim.x * gridDim.x) ????????{ ????????????row_d_matrix[i] = count; ????????????count++; ????????} ????} } intmain() { ????float*d_matrix; ????float*dc_matrix; ????dc_matrix = (float*)malloc(sizeof(float)*M*N); ????size_tpitch; ????cudaMallocPitch(&d_matrix,&pitch,M*sizeof(float),N); ????kernel<<<GridSize,BlockSize>>>(d_matrix,pitch); ????cudaMemcpy2D(dc_matrix, M * sizeof(float), d_matrix, pitch, M * sizeof(float), N, cudaMemcpyDeviceToHost); ????for(inti=0;i<M*N;i++) ????????cout<<dc_matrix[i]<<endl; ????cudaFree(d_matrix); ????free(dc_matrix); ????return0; }

上面這段代碼就是使用cudaMallocPitch來為一個N行M列的矩陣分配GPU內存空間，這里的pitch實際上就是指一行的內存大小，在這里就是M*sizeof(float)，當然有人可能會問pitch的作用是什么，一般來說，當我們使用cudaMalloc的時候，它分配的是線性內存，類似于C語言中的malloc函數，連續的內存空間，從上一個元素到訪問下一個相鄰元素的代價比較小。但是如果我們希望得到一個100*100的二維數組的話，如果我們依舊使用cudaMalloc來分配10000個內存空間的話，那我們訪問某一行的話我們就要遍歷前面的所有元素去訪問，為了減小訪問單行的代價，我們希望我們的每一行起始地址與第一行的地址是對齊的。同時，如果數組是在GPU的共享內存中，通常它會被劃分到幾個不同的bank中，這樣有多個線程訪問時就會訪問到不同的bank，如果我們希望我們的每一行可以被并行訪問的話，我門就需要保持地址對齊。cudaMallocPitch所做的事情就是：首先分配第一行的空間，并且檢查它的總字節數是否是128的倍數，如果不是的話，就再多分配幾個空余空間，使得總大小為128的倍數，這個一行的大小（包括補齊部分）就是一個pitch，然后以此類推分配其他行。最后，分配的總內存要大于實際所需的內存。因此，現在我們訪問某一行的某個元素時，就不是按照原來的想法a[row*i+j]而是使用a[pitch*i+j]來訪問，就是這個原理。因此，我們使用cudaMallocPitch的時候，一定要返回pitch的，只有這樣，我們才能訪問二維數組的某個元素。

同理，當我們需要進行二維內存復制時，假如我們希望將這個二維數組從device復制到host的時候，如果直接使用cudaMemcpy則不僅復制了數組的元素，同時也復制了補齊的內存，這是我們不希望得到的。即我們想告訴GPU我們只希望復制二維數組的元素部分，這個時候就可以使用cudaMemcpy2D來實現這個功能了，只復制有效元素，跳過補齊的內存。

總結

以上是生活随笔為你收集整理的cuda二维数组内存分配和数据拷贝的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。