從0開始學習《GPU高效能運算之CUDA》——1

文章目錄

• 函數類型
• 儲存類型
• 內建變數

0 序言

CUDA是異構編程的一個大頭，洋洋洒洒的看了寫資料，但是，感覺這個技術沒有像C++或者Java那樣有自己的權威的《編程思想》來指導系統學習，總是感覺心裡不踏實，是不是自己還沒掌握深入、或者說心裡沒底氣說自己已經入門了、已經熟悉了、已經精通了。站在一個初學者的角度，作為一個筆記式的記錄，講解自己學習和理解CUDA過程中的一些列想到的、碰到的問題。享受一個東西不一定是結果，可以是從無知到瞭解到精通的這個整個過程。

1 給自己提幾個問題

對的，我想要做什麼事情的時候，習慣性的給自己提如下問題：

問題1：GPU高效能運算之CUDA（下午簡稱CUDA技術）是幹嘛用的，我為什麼要學它？

這個問題如果我回答是，純粹為了掌握一門新的技術，如果你還是學生則可以，如果你是一個手裡有項目的工程師，那麼，我覺得沒什麼必要去學這個東西。我個人理解CUDA是電腦裡面的邊緣技術，是對程式執行效能的提高的一種方式，可以理解成一個工具。工具這種東西，你需要用的時候再去學，你不需要用的時候，你知道有這回事情就可以了。如果你用不到的東西，你也很貪心什麼都去學，第一學不精，第二電腦的技術太多太廣，對應我這樣的智商一般的人來說是不靠譜的。

我要學習CUDA，因為項目的任務可分解性較大，粒度相關性小，某項目大概一個計算任務可以分解成3000*3000*300規模的計算，而且這些元素之間完全獨立。又加入，你需要對全市100萬人，每個人計算下該個體的每年的收入與支出的淨值（從1900年計算到2000年），淡然有人說放excel中計算不就得了，好吧……我只是講個通俗點的例子。那麼這個計算規模是1000000*10，倘若要對每個人每年在做點其他什麼進階的演算法得出一個什麼指數，恐怕excel還是實現起來不太容易。所以，這種並行度很大的問題，我們可以考慮用CUDA來解決。也許你也知道，搞映像的，CUDA就顯得很重要啦。

書上說：

問題2：我的基礎是什嗎？

學習新的技術，我喜歡和之前學過的某個還熟悉的東西做為比對，這樣理解起來可能會快一些。比如學Java的時候，我想著C++，學UML的時候，我想著物件導向編程，學CUDA呢？我有什麼嗎？可能很多人會沒有異構編程的曆史。我很幸運，之前弄過半年的OpenMP，對的——多核編程技術。後來我瞭解到，如果是多個GPU，是需要用到OpenMP來做的。OpenMP本身和CUDA好像並沒什麼關係，但是，片上多核的並行演算法是很想通的。很好的一個例子是：奇偶排序。哈哈，給自己提了點學好CUDA的信心。

好了，明白自己的需求和基礎，給自己一個學習的定位，什麼地方該花時間去琢磨，相比應該很清楚。

2 你好，GPU！

GPU有兩大開發商——英偉達和AMD，支援CUDA的是英偉達，很好啊，之前買電腦是有先見之明的，買了英偉達的顯卡——GT520，不是特別高端，和單位的GTX650ti2G比起來，有那麼一點遜色，但是，好歹可以跑CUDA！

英偉達支援CUDA編程的顯卡型號從G8800開始，都是可以的。一開始作為影像處理用，而今，天文地理、數學金融、醫學軍事等等，都開始嘗試發揮GPU的優勢。

GPU的計算核心也是隔年換代，現在已經倒GK10X了，計算能力也是逐漸提升，目前已經最好的有3.0。GPU的架構從原始的到費米的，再到開普勒的，我們沒必要去一個個瞭解，我們先瞭解GPU的這個大概的曆史，免得和人交談說不出一和二，脫離菜鳥嘛！

GPU和CPU的區別可以參考下，講的還算詳細和明了。

http://www.cnblogs.com/viviman/archive/2012/11/26/2789113.html

可以這樣的去理解，單核CPU多線程並行，是感官上的並行，世界上在CPU上還是串列指令在跑；而在GPU上，才叫真正的並行！需要介紹下CUDA1.0開發包是支援在CPU上類比GPU開發的，其原理就是用多線程來類比；而現在的版本就不支援了。最新的是5.0的，官網上是下不到之前的了，反正我是找到腿軟了還沒找到。1.0是古董了！如果你有，一定要給我開開眼見哦。

我們的CUDA編程，很明顯是GPU和CPU一起來處理的嘛——異構編程！對的，我自己的理解是就一個工程而言：CPU處理串列計算業務，GPU處理並行計算業務，這裡將的並行都是並行度可觀的哦，不是說兩個元素你也來GPU上計算，那樣是不環保的——會浪費很多GPU的資源！

說道環保，我想多說一句，在GPU編程，很體現“綠色”理念。48個核，你如果寫的好，48個核全在幹活，而且乾的是有意義的活，那麼你是合理的利用資源，如果你只讓一個核在幹有意義的活，其他的都在空轉，那麼你很浪費電哦。

對於CPU和GPU分配自己的業務，稍微畫個圖失意一下，1：

圖1 工程中GPU和CPU的分工

總的來說，GPU只是幹計算並行度高的功能模組的活，一定不可以越權啦！

3 你好，CUDA！3.1 開發環境配置

第一次和同事交談，我說你和我說說C-U-D-A。他說：哭打……苦打…… ……。半天后，我說，你和我說說C-U-D-A，你說哭打是什麼東西……它說哭打就是C-U-D-A。我操，頓時傻眼了，原來這東西行業裡年哭打，對的，Cu - Da，連起來就是這麼發音的，我的無知啊。我們還是用中文解釋吧：CUDA的意思是統一計算裝置架構。

CUDA的整合式開發環境可以參考下：

http://www.cnblogs.com/viviman/archive/2012/11/05/2775100.html

在win7+vs2008+CUDA5.0的環境下體驗，是一種新的嘗試，我自己配的時候，很少或者就沒有5.0的配置部落格文檔等等。因為5.0的那一場雪比2010年來的更晚一些……

5.0是和2.0、3.0、4.0都會有那麼一點不同的，是整合了SDK和TOOL兩個東西，之前是分開的，現在是合在一起的。其實是差不多的，但是，這一分一合，就會給人很不習慣。不過，我雖然愚鈍，但是，試了幾次之後還是摸索成功。當第一個helloworld輸出後，心裡是有那麼一點小激動的，立馬跑到小區門口，買了半斤羊肉吃了，因為為了配這一套環境，我中午飯都沒吃！我這隻哭打小菜鳥就是可憐啊！

3.2 特殊的"hello world"

搭建好環境，你肯定想看看我的helloworld程式，對的，但是我不想輸出helloworld，我想幹一件事情是：我在CPU上建立個變數，傳到GPU中，然後，在GPU中賦值，然後傳出來。這件事情如果成功了，是不是可以說明，通了+GPU工作了！網上一搜的CUDA的helloworld程式，都是在CPU上輸出helloworld，那多不過癮。

__global__ void hello(char *ch)

{

ch = {'h', 'e', 'l', 'l', 'o'};

}

int main()

{

……

hello<<<1,1>>>(dev_ch);

……

return 0;

}

I think you know my idea.

在你網上搜尋到的程式的基礎上，作這樣的一個改變，相信自己動手的才是快樂的。

4 敲開編程的門

我習慣性的喜歡先看一門語言的關鍵字，CUDA的關鍵字很簡單很少：

函數類型

__global__

用來修飾核心功能的，核心功能是什麼呢，核心功能是跑在GPU上的函數；與之對應的是主機函數，用__host__修飾，也可以預設，跑在CPU上。因此，CPU也叫主機，GPU也叫裝置。通常定義這個核心功能，我喜歡在函數名前加個kernel作為修飾，讓自己清楚點。

比如__global__ void kerneladd(float *a){}

__device__

也是用來修飾核心功能的，那和__global__有什麼區別嗎？對的。__global__修飾的核心功能只能被主機函數調用；__device__修飾的核心功能只能被核心功能調用，應該很好理解。

__host__

主機函數，供主機函數調用，可預設哦，一般情況下，都是預設的，知道這個東西就行。

儲存類型

寄存器：在核函數內 int i即表示寄存器變數。

__global__：全域記憶體。在主機函數中開闢和釋放。

__shared__：共用儲存，每個block內的線程共用這個儲存。

__constant__：常量儲存，唯讀。定義在所有函數之外，作用範圍整個檔案。

__texture__：紋理儲存，唯讀。記憶體不連續。

 

內建變數

dim3

threadId

blockId

gridId