我的CUDA學習之旅——啟程

01-26

引言

本人從今年3月份開始在做學校項目（機器視覺項目）時接觸到了並行編程這一概念，項目的主要目的是對蠶繭的雌雄以及品質進行檢測以及分選，那時候的目的是為了在圖像識別的時候通過多個線程同時對多張傳入的圖片進行並行處理以達到加速程序運行速度，運用的方法主要是利用了C++自帶的future庫以及thread庫，具體實現方法可參照此博客[C++11 並髮指南四](http://www.cnblogs.com/haippy/p/3239248.html)

此方法簡單來說是利用了目前計算機為多核式結構（四核），因為在一般的串列編程中，一個程序的指令在單一的 CPU 上按照先後順序依次執行，而並行編程則將一個程序分成獨立的若干部分在一個或多個 CPU 上進行同步執行以取得更高的運算效率和性能。在我自己的項目中我將圖像檢測的演算法寫入一個函數中（演算法實現較為簡單），通過上述方法，同時對3張傳入的待檢測圖進行識別（攝像機拍攝到的蠶繭圖片），部分代碼如下

Mat img1 = imread("new.bmp");//讀入第一張圖片 Mat img2 = imread("new1.bmp");//讀入第二張圖片 Mat img3 = imread("new2.bmp");//讀入第三張圖片 Mat imgResult1, imgResult2, imgResult3;//創建識別結果 //多線程運算 auto f1 = async(Pod, img1);//Pod為檢測演算法函數 auto f2 = async(Pod, img2); auto f3 = async(Pod, img3); imgResult1 = f1.get(); imgResult2 = f2.get(); imgResult3 = f3.get();

但由於CPU加速效果的有限，在處理一些大數據的時候，這種並行的方式並不能從本質上解決程度處理速度慢的問題，故在5月份有幸接觸到了GPU並行加速這一概念，開始學習Nvidia公司推出的基於CUDA的並行程序設計，到目前本人已經完成了一些基於OPENCV的圖像處理演算法關於CUDA的改編，在加速圖像處理程序上有著不錯的效果，故打算將自己的一些關於CUDA的學習以及實踐進行記錄，如有不當之處，希望各位大佬們批評指正~

CUDA編程基礎

了解電腦的同學都知道現在每台計算機在其硬體構成上都包括最重要的兩部分——主機和顯卡，而在CUDA的架構中引入了主機端（host）和設備（device）的概念，主機端即為CPU，而設備端則為GPU，所謂的CUDA編程即為利用GPU在大規模並行計算上的執行效率高的優勢（一些對顯卡要求高的大型遊戲本質上來講就是要求顯卡的計算速度快），利用CPU，將一些邏輯複雜性較小的計算部分分配給GPU讓其協助完成計算任務。

一個完整的CUDA程序由主機代碼和設備代碼兩部分組成。主機端代碼部分在CPU上執行，是普通的C代碼；設備端代碼部分在GPU上執行，此代碼部分在內核kernel上遍寫（.cu文件）。關於CUDA中C語言的編寫規範以及一些功能網上已有許多健全的資料，在這裡不在多作解釋，簡單利用一個CUDA程序進行作為常式進行實踐：

#include "cuda_runtime.h"#include "device_launch_parameters.h"#include <stdlib.h>#include <iostream>using namespace std;//將數組A與B中的元素相加存入數組C中__global__ void arrayAdd(int *A, int *B, int *C){ int i = threadIdx.x; C[i] = A[i] + B[i];}int main(){ //簡單的示例，A、B、C均為長度為4的整型數組 //定義主機端數組A、B、C int A[4] = { 1, 2, 3, 4 }; int B[4] = { 2, 3, 4, 5 }; int C[4]; //定義設備端數組d_a、d_b、d_c int *d_a; int *d_b; int *d_c; //定義他們在設備端的空間大小 int size = 4 * sizeof(int); //在GPU中為他們開闢分配對應的顯存空間 cudaMalloc((void**)&d_a, size); cudaMalloc((void**)&d_b, size); cudaMalloc((void**)&d_c, size); //利用cudaMemcpy函數在CPU端A,B的值複製到對應的GPU內存中 cudaMemcpy(d_a, A, size, cudaMemcpyHostToDevice); cudaMemcpy(d_b, B, size, cudaMemcpyHostToDevice); //調用編寫好的內核程序，實現數組相加功能 arrayAdd << <1, 4 >> >(d_a, d_b, d_c); //利用cudaMemcpy函數將GPU端計算結果複製到CPU端 cudaMemcpy(C, d_c, size, cudaMemcpyDeviceToHost); //釋放所有的GPU內存（切記有借有還，不然後果自負。。。） cudaFree(d_a); cudaFree(d_b); cudaFree(d_c); system("Pause"); return 0;}

CUDA編程簡單來說分為3大步，首先創建在GPU中的內存空間，將有初值的變數從CPU端複製值GPU端，接著調用在GPU內運行的內核函數進行計算，最後將內核函數的運算結果複製到CPU端的變數並且釋放掉所有之前創建的GPU內存（這步不能漏，不然後面會出一堆內存問題）。

以上面的程序為例，程序目的為利用並行計算將數組A與數組B相加，首先我在主機端隨意定義了3個int型數組，其數組長度均為4，接著利用cudaMalloc函數對其GPU內存進行開闢，內存大小為4倍的int型，接著利用cudaMemcpy函數對數據進行傳輸，這裡在函數的末尾「cudaMemcpyHostToDevice」為一個cudaMemcpy函數特有的標識碼，即將主機端（Host）的數據傳入設備端（Device）。

然後調用了內核函數arrayADD，在解釋此內核函數前，首先我將引入CUDA中Grid（網格）、Block（塊）、Thread（線程）的概念，每個內核函數在啟動前，程序員必須對本次計算的維度進行設計，thread是最小的計算單元，多個thread組成了一個block,而多個block組成了一個grid，這裡可能有同學會問，讓多線程一起計算直接用相應數量的thread不就行了，為什麼還要引入block和grid，增加設計的複雜度，這是因為每台電腦的顯卡在硬體設計的時候都存在的一定的物理極限，簡單來說一個block中的thread不是無限的，舉個例子，本人目前使用的本本中，顯卡型號為Nvidia Geforce 755M，它的最大線程數為1024，超過這個值程序就無法運行，故在大計算時需要合理設計block中thread的數量，這在後續的實踐中我會繼續提到，關於Grid（網格）、Block（塊）、Thread（線程）這三者更具體的關係可參見此博客[CUDA 的Threading：Block 和Grid 的設定與Warp](http://www.360doc.com/content/13/0819/19/12278894_308350397.shtml)

由於本例數據較少，在計算任務的分配上直接利用了4個thread和一個block，如此設計在於每個thread承擔一次A[i]+B[i]的任務。最後再利用cudaMemcpy函數將運算結果d_c傳送給數組C，一個簡單的CUDA小程序就完成了，由於本例的目的在於初步嘗試CUDA編程，一些細節小問題望大家不要深究~

CUDA中線程索引threadIdx和塊索引blockIdx

在上文的內核函數中，第一行中threadIdx.x為線程的索引，它的目的在於讓GPU可以準確的找到目的操作數用對應的線程進行計算工作，是CUDA編程中和重要的一部分。

本人在學習CUDA的過程中，一開始經常被線程索引threadIdx和塊索引blockIdx搞的崩潰，現將我在學習中總結一些的小經驗進行分享：

（1）CUDA中kernel內的變數和常量

在CUDA中每個塊中線程的索引變數threadIdx分為X方向上的索引threadIdx.x以及Y方向上的索引threadIdx.y，塊索引變數blockIdx分為X方向上的索引blockIdx.x以及Y方向上的索引block.y，通過這兩個索引我們可以準確定位到某個線程。

而kernel中的常量為blockDim以及gridDim，它們分別代表了沒block的大小以及grid中block的分布。

（2）創建內核程序中block和thread

dim3 blocksPerGrid(N1,M1);//在grid中有N1*M1個blockdim3 threadsPerBlock(N2,M2);//每個block中有N2*M2個thread

（3）創建索引

int threadIndex = threadIdx.x + threadIdx.y * blockDim.x;int blockIndex = blockIdx.x + blockIdx.y * gridDim.x;int index = threadIndex + blockIndex * blockDim.x * blockDim.y;

threadIndex為單個block中thread的索引，blockIndex為block在整個grid中的索引，index即為某個thread在整個grid中的索引。

結束

這是我的第一篇CUDA知乎文，寫的可能有些粗陋，後續會繼續上傳我的一些CUDA程序，希望大家多多指教~