CUDA: GPU高性能运算
CUDA: GPU高性能運算
2013-10-11 22:23 5650人閱讀 評論(0) 收藏 舉報 分類： CUDA（106）目錄(?)[+]
0?序言
CUDA是異構(gòu)編程的一個大頭，洋洋灑灑的看了些資料，但是，感覺這個技術(shù)沒有像C++或者Java那樣有自己的權(quán)威的《編程思想》來指導(dǎo)系統(tǒng)學(xué)習(xí)，總是感覺心里不踏實，是不是自己還沒掌握深入、或者說心里沒底氣說自己已經(jīng)入門了、已經(jīng)熟悉了、已經(jīng)精通了。站在一個初學(xué)者的角度，作為一個筆記式的記錄，講解自己學(xué)習(xí)和理解CUDA過程中的一些列想到的、碰到的問題。享受一個東西不一定是結(jié)果，可以是從無知到了解到精通的這個整個過程。
1?給自己提幾個問題
對的，我想要做什么事情的時候，習(xí)慣性的給自己提如下問題：
問題1：GPU高性能運算之CUDA（下午簡稱CUDA技術(shù)）是干嘛用的，我為什么要學(xué)它？
這個問題如果我回答是，純粹為了掌握一門新的技術(shù)，如果你還是學(xué)生則可以，如果你是一個手里有項目的工程師，那么，我覺得沒什么必要去學(xué)這個東西。我個人理解CUDA是計算機里面的邊緣技術(shù)，是對程序執(zhí)行性能的提高的一種方式，可以理解成一個工具。工具這種東西，你需要用的時候再去學(xué)，你不需要用的時候，你知道有這回事情就可以了。如果你用不到的東西，你也很貪心什么都去學(xué)，第一學(xué)不精，第二計算機的技術(shù)太多太廣，對應(yīng)我這樣的智商一般的人來說是不靠譜的。
我要學(xué)習(xí)CUDA，因為項目的任務(wù)可分解性較大，粒度相關(guān)性小，某項目大概一個計算任務(wù)可以分解成3000*3000*300規(guī)模的計算，而且這些元素之間完全獨立。又加入，你需要對全市100萬人，每個人計算下該個體的每年的收入與支出的凈值（從1900年計算到2000年），淡然有人說放excel中計算不就得了，好吧……我只是講個通俗點的例子。那么這個計算規(guī)模是1000000*10，倘若要對每個人每年在做點其他什么高級的算法得出一個什么指數(shù)，恐怕excel還是實現(xiàn)起來不太容易。所以，這種并行度很大的問題，我們可以考慮用CUDA來解決。也許你也知道，搞圖像的，CUDA就顯得很重要啦。
問題2：我的基礎(chǔ)是什么？
學(xué)習(xí)新的技術(shù)，我喜歡和之前學(xué)過的某個還熟悉的東西做為比對，這樣理解起來可能會快一些。比如學(xué)Java的時候，我想著C++，學(xué)UML的時候，我想著面向?qū)ο缶幊?#xff0c;學(xué)CUDA呢？我有什么么？可能很多人會沒有異構(gòu)編程的歷史。我很幸運，之前弄過半年的OpenMP，對的——多核編程技術(shù)。后來我了解到，如果是多個GPU，是需要用到OpenMP來做的。OpenMP本身和CUDA好像并沒什么關(guān)系，但是，片上多核的并行算法是很想通的。很好的一個例子是：奇偶排序。哈哈，給自己提了點學(xué)好CUDA的信心。
好了，明白自己的需求和基礎(chǔ)，給自己一個學(xué)習(xí)的定位，什么地方該花時間去琢磨，相比應(yīng)該很清楚。
2?你好，GPU！
GPU有兩大開發(fā)商——英偉達和AMD，支持CUDA的是英偉達，很好啊，之前買電腦是有先見之明的，買了英偉達的顯卡——GT520，不是特別高端，和單位的GTX650ti2G比起來，有那么一點遜色，但是，好歹可以跑CUDA！
英偉達支持CUDA編程的顯卡型號從G8800開始，都是可以的。一開始作為圖像處理用，而今，天文地理、數(shù)學(xué)金融、醫(yī)療軍事等等，都開始嘗試發(fā)揮GPU的優(yōu)勢。
GPU的計算核心也是隔年換代，現(xiàn)在已經(jīng)倒GK10X了，計算能力也是逐漸提升，目前已經(jīng)最好的有3.0。GPU的架構(gòu)從原始的到費米的，再到開普勒的，我們沒必要去一個個了解，我們先了解GPU的這個大概的歷史，免得和人交談?wù)f不出一和二，脫離菜鳥嘛！
GPU和CPU的區(qū)別可以參考下，講的還算詳細和明了。
http://www.cnblogs.com/viviman/archive/2012/11/26/2789113.html
可以這樣的去理解，單核CPU多線程并行，是感官上的并行，世界上在CPU上還是串行指令在跑；而在GPU上，才叫真正的并行！需要介紹下CUDA1.0開發(fā)包是支持在CPU上模擬GPU開發(fā)的，其原理就是用多線程來模擬；而現(xiàn)在的版本就不支持了。最新的是5.0的，官網(wǎng)上是下不到之前的了，反正我是找到腿軟了還沒找到。1.0是古董了！如果你有，一定要給我開開眼見哦。
我們的CUDA編程，很明顯是GPU和CPU一起來處理的嘛——異構(gòu)編程！對的，我自己的理解是就一個工程而言：CPU處理串行計算業(yè)務(wù)，GPU處理并行計算業(yè)務(wù)，這里將的并行都是并行度可觀的哦，不是說兩個元素你也來GPU上計算，那樣是不環(huán)保的——會浪費很多GPU的資源！
說道環(huán)保，我想多說一句，在GPU編程，很體現(xiàn)“綠色”理念。48個核，你如果寫的好，48個核全在干活，而且干的是有意義的活，那么你是合理的利用資源，如果你只讓一個核在干有意義的活，其他的都在空轉(zhuǎn)，那么你很浪費電哦。
對于CPU和GPU分配自己的業(yè)務(wù)，稍微畫個圖失意一下，如圖1：
圖1?工程中GPU和CPU的分工
總的來說，GPU只是干計算并行度高的功能模塊的活，一定不可以越權(quán)啦！
3?你好，CUDA！
3.1?開發(fā)環(huán)境配置
第一次和同事交談，我說你和我說說C-U-D-A。他說：哭打……苦打……?……。半天后，我說，你和我說說C-U-D-A，你說哭打是什么東西……它說哭打就是C-U-D-A。我操，頓時傻眼了，原來這東西行業(yè)里年哭打，對的，Cu?-?Da，連起來就是這么發(fā)音的，我的無知啊。我們還是用中文解釋吧：CUDA的意思是統(tǒng)一計算設(shè)備架構(gòu)。
CUDA的集成開發(fā)環(huán)境可以參考下：
http://www.cnblogs.com/viviman/archive/2012/11/05/2775100.html
在win7+vs2008+CUDA5.0的環(huán)境下體驗，是一種新的嘗試，我自己配的時候，很少或者就沒有5.0的配置博客文檔等等。因為5.0的那一場雪比2010年來的更晚一些……
5.0是和2.0、3.0、4.0都會有那么一點不同的，是集成了SDK和TOOL兩個東西，之前是分開的，現(xiàn)在是合在一起的。其實是差不多的，但是，這一分一合，就會給人很不習(xí)慣。不過，我雖然愚鈍，但是，試了幾次之后還是摸索成功。當?shù)谝粋€helloworld輸出后，心里是有那么一點小激動的，立馬跑到小區(qū)門口，買了半斤羊肉吃了，因為為了配這一套環(huán)境，我中午飯都沒吃！我這只哭打小菜鳥就是可憐啊！
3.2?特殊的"hello?world"
搭建好環(huán)境，你肯定想看看我的helloworld程序，對的，但是我不想輸出helloworld，我想干一件事情是：我在CPU上創(chuàng)建個變量，傳到GPU中，然后，在GPU中賦值，然后傳出來。這件事情如果成功了，是不是可以說明，通了+GPU工作了！網(wǎng)上一搜的CUDA的helloworld程序，都是在CPU上輸出helloworld，那多不過癮。
[cpp] view plaincopyI?think?you?know?my?idea.
在你網(wǎng)上搜索到的程序的基礎(chǔ)上，作這樣的一個改變，相信自己動手的才是快樂的。
4?敲開編程的門
我習(xí)慣性的喜歡先看一門語言的關(guān)鍵字，CUDA的關(guān)鍵字很簡單很少：
函數(shù)類型
__global__:?用來修飾內(nèi)核函數(shù)的，內(nèi)核函數(shù)是什么呢，內(nèi)核函數(shù)是跑在GPU上的函數(shù)；與之對應(yīng)的是主機函數(shù)，用__host__修飾，也可以缺省，跑在CPU上。因此，CPU也叫主機，GPU也叫設(shè)備。通常定義這個內(nèi)核函數(shù)，我喜歡在函數(shù)名前加個kernel作為修飾，讓自己清楚點。
比如__global__?void?kerneladd(float?*a){}
__device__: 也是用來修飾內(nèi)核函數(shù)的，那和__global__有什么區(qū)別嗎？對的。__global__修飾的內(nèi)核函數(shù)只能被主機函數(shù)調(diào)用；__device__修飾的內(nèi)核函數(shù)只能被內(nèi)核函數(shù)調(diào)用，應(yīng)該很好理解。
__host__: 主機函數(shù)，供主機函數(shù)調(diào)用，可缺省哦，一般情況下，都是缺省的，知道這個東西就行。
存儲類型
寄存器：在核函數(shù)內(nèi)?int?i即表示寄存器變量。
__global__： ? ? 全局內(nèi)存。在主機函數(shù)中開辟和釋放。
__shared__： ? 共享存儲，每個block內(nèi)的線程共享這個存儲。
__constant__：常量存儲，只讀。定義在所有函數(shù)之外，作用范圍整個文件。
__texture__： ? 紋理存儲，只讀。內(nèi)存不連續(xù)。
內(nèi)建變量
dim3, threadId, blockId, ?gridId
5?GPU也不允許偏心
并行的事情多了，我們作為GPU的指令分配者，不能偏心了——給甲做的事情多，而乙沒事做，個么甲肯定不爽的來。所以，在GPU中，叫做線程網(wǎng)絡(luò)的分配。首先還是來看下GPU的線程網(wǎng)絡(luò)吧，圖2：
圖2?線程網(wǎng)絡(luò)
我們將具體點的，在主機函數(shù)中如果我們分配的是這樣的一個東西：
dim3?blocks(32, 32);
dim3?threads(16, 16);
dim3是神馬？dim3是一個內(nèi)置的結(jié)構(gòu)體，和linux下定義的線程結(jié)構(gòu)體是個類似的意義的東西，dim3結(jié)構(gòu)變量有x，y，z，表示3維的維度。不理解沒關(guān)系，慢慢看。
kernelfun<<<blocks,?threads>>>();
我們調(diào)用kernelfun這個內(nèi)核函數(shù)，將blocks和threads傳到<<<,>>>里去，這句話可牛逼大了——相當于發(fā)號施令，命令那些線程去干活。這里使用了32*32?*?16*16個線程來干活。你看明白了嗎？blocks表示用了二維的32*32個block組，而每個block中又用了16*16的二維的thread組。好吧，我們這個施令動用了262144個線程！我們先不管GPU內(nèi)部是如何調(diào)度這些線程的，反正我們這一句話就是用了這么多線程。
那我們的內(nèi)核函數(shù)kernelfun()如何知道自己執(zhí)行的是哪個線程？這就是線程網(wǎng)絡(luò)的特點啦，為什么叫網(wǎng)絡(luò)，是有講究的，網(wǎng)絡(luò)就可以定格到網(wǎng)點：
比如int?tid?=?threadId.x?+?blockId.x?*?16
這里有一個講究，block是有維度的，一維、二維、三維。
對于一維的block，tid?=?threadId.x
對于（Dx，Dy）二維的block，tid?=?threadId.x?+?Dx*threadId.y
對于（Dx，Dy，Dz）三維的block，tid?=?threadId.x?+?Dx*threadId.y?+?Dz*Dy*threadId.z
我習(xí)慣的用這樣的模式去分配，比較通用：
dim3?dimGrid();
dim3?dimBlock();
kerneladd<<<dimGrid,?dimBlock>>>();
這可是萬金油啊，你需要做的事情是填充dimGrid和dimBlock的結(jié)構(gòu)體構(gòu)造函數(shù)變量，比如，dimGrid(16,?16)表示用了16*16的二維的block線程塊。
（0,0）（0,1）（0,2）……（0,15）
（1,0）（1,1）（1,2）……（1,15）
（2,0）（2,1）（2,2）……（2,15）
……
（15,0）（15,1）（15,2）……（15,15）
（，）是（dimGrid.x，?dimGrid.y）的網(wǎng)格編號。
我們這么理解吧，現(xiàn)在又一群人，我們分成16*16個小組（block），排列好，比如第3行第4列就指的是（2,3）這個小組。
而dimBlock（16,16）表示每個小組有16*16個成員，如果你想點名第3行第4列這個小組的里面的第3行第4列那個同學(xué)，那么，你就是在（2,3）這個block中選擇了（2,3）這個線程。這樣應(yīng)該有那么一點可以理解進去的意思了吧？不理解透徹么什么關(guān)系，這個東西本來就是cuda中最讓我糾結(jié)的事情。我們且不管如何分配線程，能達到最優(yōu)化，我們的目標是先讓GPU正確地跑起來，計算出結(jié)果即可，管他高效不高效，管他環(huán)保不環(huán)保。
嘮叨了這么多，下面我們用一個最能說明問題的例子來進一步理解線程網(wǎng)絡(luò)分配機制來了解線程網(wǎng)絡(luò)的使用。
一維網(wǎng)絡(luò)線程
eg：int?arr[1000]，對每個數(shù)組元素進行加1操作。
idea：我們最直接的想法，是調(diào)度1000個線程去干這件事情。
first?pro：我想用一個小組的1000個人員去干活。這里會存在這樣一個問題——一個小組是不是有這么多人員呢？是的，這個事情你必須了解，連自己組內(nèi)多少人都不知道，你也不配作指揮官呀。對的，這個參數(shù)叫做maxThreadsPerBlock，如何取得呢？
好吧，cuda定義了一個結(jié)構(gòu)體cudaDeviceProp，里面存入了一系列的結(jié)構(gòu)體變量作為GPU的參數(shù)，出了maxThreadsPerBlock，還有很多信息哦，我們用到了再說。
maxThreadsPerBlock這個參數(shù)值是隨著GPU級別有遞增的，早起的顯卡可能512個線程，我的GT520可以跑1024個線程，辦公室的GTX650ti2G可以跑1536個，無可非議，當然多多益善。一開始，我在想，是不是程序?qū)⒚總€block開的線程開滿是最好的呢？這個問題留在以后在說，一口吃不成胖子啦。
好吧，我們的數(shù)組元素1000個，是可以在一個block中干完的。
內(nèi)核函數(shù)：
[cpp] view plaincopy
呀，原來這么簡單，個么CUDA也忒簡單了哇！這中想法是好的，給自己提高信心，但是這種想法多了是不好的，因為后面的問題多了去了。
盆友說，1000個元素，還不如CPU來的快，對的，很多情況下，數(shù)據(jù)量并行度不是特別大的情況下，可能CPU來的更快一些，比較設(shè)備與主機之間互相調(diào)度操作，是會有額外開銷的。有人就問了，一個10000個元素的數(shù)組是不是上面提供的idea就解決不了啦？對，一個block人都沒怎么多，如何完成！這個情況下有兩條路可以選擇——
第一，我就用一個組的1000人來干活話，每個人讓他干10個元素好了。
這個解決方案，我們需要修改的是內(nèi)核函數(shù)：
[cpp] view plaincopy第二，我多用幾個組來干這件事情，比如我用10個組，每個組用1000人。
這個解決方案就稍微復(fù)雜了一點，注意只是一點點哦~因為，組內(nèi)部怎么干活和最原始的做法是一樣的，不同之處是，我們調(diào)遣了10個組去干這件事情。
首先我們來修改我們的主機函數(shù)：
[cpp] view plaincopy盆友要問了，10個組每個組1000人，你怎么點兵呢？很簡單啊，第1組第3個線程出列，第9組第9個線程出列。每個人用組號和組內(nèi)的編號定了位置。在線程網(wǎng)絡(luò)中，blockId.x和threadId.x就是對應(yīng)的組號和組內(nèi)編號啦，我必須要這里開始形象點表示這個對應(yīng)關(guān)系，如果這個對應(yīng)關(guān)系是這樣子的[blockId.x，threadId.x]，那么我們的數(shù)組arr[10000]可以這樣分配給這10個組去干活：
(0,0)——arr[0]，(0,1)——arr[1]，……(0,999)——arr[999]
(1,0)——arr[0+1*1000]，(1,1)——arr[1+1*1000]，……(1,999)——arr[999+1*1000]
……
(9,0)——arr[0+9*1000]，(9,1)——arr[1+9*1000]，……(9,999)——arr[999+9*1000]
是不是很有規(guī)律呢？對的，用blockId.x和threadId.x可以很好的知道哪個線程干哪個元素，這個元素的下表就是threadId.x?+?1000*blockId.x。
這里我想說的是，如果我們哪天糊涂了，畫一畫這個對應(yīng)關(guān)系的表，也許，就更加清楚的知道我們分配的線程對應(yīng)的處理那些東西啦。?
一維線程網(wǎng)絡(luò)，就先學(xué)這么多了。
二維網(wǎng)絡(luò)線程
eg2：int?arr[32][16]二維的數(shù)組自增1。
第一個念頭，開個32*16個線程好了哇，萬事大吉！好吧。但是，朕現(xiàn)在想用二維線程網(wǎng)絡(luò)來解決，因為朕覺得一個二維的網(wǎng)絡(luò)去映射一個二維的數(shù)組，朕看的更加明了，看不清楚自己的士兵，如何帶兵打仗！
我還是畫個映射關(guān)系：
一個block中，現(xiàn)在是一個二維的thread網(wǎng)絡(luò)，如果我用了16*16個線程。
(0,0)，(0,1)，……(0,15)
(1,0)，(1,1)，……(1,15)
……
(15,0)，(15,1)，……(15,15)
呀，現(xiàn)在一個組內(nèi)的人稱呼變了嘛，一維網(wǎng)絡(luò)中，你走到一個小組里，叫3號出列，就出來一個，你現(xiàn)在只是叫3號，沒人會出來！這個場景是這樣的，現(xiàn)在你班上有兩個人同名的人，你只叫名，他們不知道叫誰，你必須叫完整點，把他們的姓也叫出來。所以，二維網(wǎng)絡(luò)中的(0,3)就是原來一維網(wǎng)絡(luò)中的3，二維中的(i,j)就是一維中的(j+i*16)。不管怎么樣，一個block里面能處理的線程數(shù)量總和還是不變的。
一個grid中，block也可以是二維的，一個block中已經(jīng)用了16*16的thread了，那我們一共就32*16個元素，我們用2個block就行了。
先給出一個代碼清單吧，程序員都喜歡看代碼，這段代碼是我抄襲的。第一次這么完整的放上代碼，因為我覺得這個代碼可以讓我說明我想說的幾個問題：
第一，二維數(shù)組和二維指針的聯(lián)系。
第二，二維線程網(wǎng)絡(luò)。
第三，cuda的一些內(nèi)存操作，和返回值的判斷。
[cpp] view plaincopy簡要的來學(xué)習(xí)一下二維網(wǎng)絡(luò)這個知識點，?
dim3?dimBlock(16, 16);?
//定義block內(nèi)的thread二維網(wǎng)絡(luò)為16*16
dim3?dimGrid((COLS + dimBlock.x - 1) / (dimBlock.x), ?(ROWS + dimBlock.y - 1) / (dimBlock.y));?
//定義grid內(nèi)的block二維網(wǎng)絡(luò)為1*2
unsigned?int?row?=?blockDim.y * blockIdx.y ?+ ?threadIdx.y;?
//二維數(shù)組中的行號
unsigned?int?col?=?blockDim.x * blockIdx.x ?+ ?threadIdx.x;?
//二維線程中的列號
三維網(wǎng)絡(luò)線程
dim3定義了三維的結(jié)構(gòu)，但是，貌似二維之內(nèi)就能處理很多事情啦，所以，我放棄學(xué)習(xí)三維。網(wǎng)上看到的不支持三維網(wǎng)絡(luò)是什么意思呢？先放一放。
給自己充充電
同一塊顯卡，不管你是二維和三維或一維，其計算能力是固定的。比如一個block能處理1024個線程，那么，一維和二維線程網(wǎng)絡(luò)是不是處理的線程數(shù)一樣呢？
?
回答此問題，先給出網(wǎng)絡(luò)配置的參數(shù)形式——<<<Dg,Db,Ns,S>>>，各個參數(shù)含義如下：
Dg：定義整個grid的維度，類型Dim3，但是實際上目前顯卡支持兩個維度，所以，dim3<<Dg.x,?Dg.y,?1>>>第z維度默認只能為1，上面顯示出這個最大有65536*65536*1，每行有65536個block，每列有65536個block，整個grid中一共有65536*65536*1個block。
Db：定義了每個block的維度，類型Dim3，比如512*512*64，這個可以定義3維尺寸，但是，這個地方是有講究了，三個維度的積是有上限的，對于計算能力1.0、1.1的GPU，這個值不能大于768，對于1.2、1.3的不能大于1024，對于我們試一試的這塊級別高點的，不能大于1536。這個值可以獲取哦——maxThreadsPerBlock
Ns：這個是可選參數(shù)，設(shè)定最多能動態(tài)分配的共享內(nèi)存大小，比如16k，單不需要是，這個值可以省略或?qū)?/span>0。
S：也是可選參數(shù)，表示流號，默認為0。流這個概念我們這里不說。
接著，我想解決幾個你肯定想問的兩個問題，因為我看很多人想我這樣的問這個問題：
1?block內(nèi)的thread我們是都飽和使用嗎？
答：不要，一般來說，我們開128或256個線程，二維的話就是16*16。
2?grid內(nèi)一般用幾個block呢？
答：牛人告訴我，一般來說是你的流處理器的4倍以上，這樣效率最高。
回答這兩個問題的解釋，我想抄襲牛人的一段解釋，解釋的好的東西就要推廣呀：
GPU的計算核心是以一定數(shù)量的Streaming?Processor(SP)組成的處理器陣列，NV稱之為Texture?Processing?Clusters(TPC)，每個TPC中又包含一定數(shù)量的Streaming?Multi-Processor(SM)，每個SM包含8個SP。SP的主要結(jié)構(gòu)為一個ALU（邏輯運算單元），一個FPU（浮點運算單元）以及一個Register?File(寄存器堆)。SM內(nèi)包含有一個Instruction?Unit、一個Constant?Memory、一個Texture?Memory，8192個Register、一個16KB的Share?Memory、8個Stream?Processor(SP)和兩個Special?Function?Units（SFU）。（GeForce9300M?GS只擁有1個SM） Thread是CUDA模型中最基本的運行單元，執(zhí)行最基本的程序指令。Block是一組協(xié)作Thread，Block內(nèi)部允許共享存儲，每個Block最多包含512個Thread。Grid是一組Block，共享全局內(nèi)存。Kernel是在GPU上執(zhí)行的核心程序，每一個Grid對應(yīng)一個Kernel任務(wù)。 在程序運行的時候，實際上每32個Thread組成一個Warp，每個?warp?塊都包含連續(xù)的線程，遞增線程?ID?。Warp是MP的基本調(diào)度單位，每次運行的時候，由于MP數(shù)量不同，所以一個Block內(nèi)的所有Thread不一定全部同時運行，但是每個Warp內(nèi)的所有Thread一定同時運行。因此，我們在定義Block?Size的時候應(yīng)使其為Warp?Size的整數(shù)倍，也就是Block?Size應(yīng)為32的整數(shù)倍。理論上Thread越多，就越能彌補單個Thread讀取數(shù)據(jù)的latency?，但是當Thread越多，每個Thread可用的寄存器也就越少，嚴重的時候甚至能造成Kernel無法啟動。因此每個Block最少應(yīng)包含64個Thread，一般選擇128或者256，具體視MP數(shù)目而定。一個MP最多可以同時運行768個Thread，但每個MP最多包含8個Block，因此要保持100%利用率，Block數(shù)目與其Size有如下幾種設(shè)定方式： ??2?blocks?x?384?threads ??3?blocks?x?256?threads ??4?blocks?x?192?threads ??6?blocks?x?128?threads ??8?blocks?x?96?threads?
這些電很重要啊，必須要充！不然，我就很難理解為什么網(wǎng)絡(luò)線程如何分配的。
6?規(guī)約思想和同步概念
擴大點說，并行計算是有一種基本思想的，這個算法能解決很多很常規(guī)的問題，而且很實用，比如說累加和累積等——規(guī)約思想。對于基礎(chǔ)的、重要的，我想有必要系統(tǒng)的學(xué)習(xí)。
我覺得有必要重新復(fù)制下之前寫的這篇介紹：
http://www.cnblogs.com/viviman/archive/2012/11/21/2780286.html
并行程序的開發(fā)有其不同于單核程序的特殊性，算法是重中之重。根據(jù)不同業(yè)務(wù)設(shè)計出不同的并行算法，直接影響到程序的效率。因此，如何設(shè)計并行程序的算法，似乎成為并編程的最大難點。觀其算法，包括cuda?sdk的例子和網(wǎng)上的牛人，給出的一些例子，以矩陣和矢量處理為主，深入點的包括fft和julia等數(shù)學(xué)公式，再高級一點的算是圖形處理方面的例子。學(xué)習(xí)這些算法的思想，免不了有自己的一點點總結(jié)。之前學(xué)習(xí)過omp編程，結(jié)合現(xiàn)在的cuda，我覺得要理解并行編程，首先理解劃分和規(guī)約這兩個概念。也許你的算法學(xué)的更加扎實。劃分是《算法》里面的一個重要思想，將一個大的問題或任務(wù)，分解成小問題小任務(wù)，各個擊破，最后歸并結(jié)果；規(guī)約是《cuda**》書上介紹的一個入門的重要思想，規(guī)約算法(reduction)用來求連加、連乘、最值等，應(yīng)用廣泛。每次循環(huán)參加運算的線程減少一半。不管算法的思想如何花樣，萬變不離其中的一點--將一個大的任務(wù)分解成小的任務(wù)集合，分解原則是粒度合適盡量小、數(shù)據(jù)相關(guān)性盡量小。如此而已。因為，我們用GPU是為了加速，要加速必須提高執(zhí)行任務(wù)的并行度！明白這個道理，那么我們將絞盡腦汁地去想方設(shè)法分析自己手上的任務(wù)，分解、分解、分解！這里拿規(guī)約來說事情，因為，規(guī)約這個東西，似乎可以拿來單做9*9乘法表來熟悉，熟悉了基礎(chǔ)的口訣，那么99*99的難題也會迎刃而解。ex：矢量加法，要實現(xiàn)N=64*256長度的矢量的累加和。假設(shè)a+b計算一次耗時t。
cpu計算：顯然單核的話需要64*256*t。我們?nèi)萑滩涣恕?/span>
gpu計算：最初的設(shè)想，我們?nèi)绻袀€gpu能同時跑N/2個線程，我們這N/2個線程同時跑，那么不是只需要t時間就能將N個數(shù)相加編程N/2個數(shù)相加了嗎？對的。這一輪我們用了t時間；接著的想法，我們不斷的遞歸這個過程，能發(fā)現(xiàn)嗎？第二輪，我們用N/2/2個線程同時跑，剩下N/2/2個數(shù)相加，這一輪我們同樣用了t時間；一直這樣想下去吧，最后一輪，我們用了1個線程跑，剩下1個數(shù)啦，這就是我們的結(jié)果！每一輪時間都為t，那么理想情況，我們用了多少輪這樣的計算呢？計算次數(shù)=log(N)=6*8=48，對的，只用了48輪，也就是說，我們花了48*t的時間！
規(guī)約就是這樣，很簡單，很好用，我們且不管程序后期的優(yōu)化，單從這個算法分析上來說，從時間復(fù)雜度N降到了logN，這在常規(guī)算法上，要提高成這樣的效率，是不得了的，這是指數(shù)級別的效率提高！所以，你會發(fā)現(xiàn)，GPU有CPU無法取代的得天獨厚的優(yōu)勢——處理單元真心多啊！
規(guī)約求和的核函數(shù)代碼如下：
[cpp] view plaincopy這個例子還讓我學(xué)到另一個東西——同步！我先不說同步是什么，你聽我說個故事：我們調(diào)遣了10個小組從南京去日本打仗，我們的約定是，10個組可以自己行動，所有組在第三天在上海機場會合，然后一起去日本。這件事情肯定是需要處理的，不能第1組到了上海就先去日本了，這些先到的組，唯一可以做的事情是——等待！這個先來后到的事情，需要統(tǒng)一管理的時候，必須同步一下，在上海這個地方，大家統(tǒng)一下步調(diào)，快的組等等慢的組，然后一起干接下去的旅程。
是不是很好理解，這就是同步在生活中的例子，應(yīng)該這樣說，計算機的所有機制和算法很多都是源于生活！結(jié)合起來，理解起來會簡單一點。
在CUDA中，我們的同步機制用處大嗎？又是如何用的呢？我告訴你，一個正常規(guī)模的工程中，一般來說數(shù)據(jù)都會有先來后到的關(guān)系，這一個計算結(jié)果可能是提供給另一個線程用的，這種依賴關(guān)系存在，會造成同步的應(yīng)用。
__synctheads()這句話的作用是，這個block中的所有線程都執(zhí)行到此的時候，都聽下來，等所有都執(zhí)行到這個地方的時候，再往下執(zhí)行。
7?撬開編程的鎖
鎖是數(shù)據(jù)相關(guān)性里面肯定要用到的東西，很好，生活中也一樣，沒鎖，家里不安全；GPU中沒鎖，數(shù)據(jù)會被“盜”。
對于存在競爭的數(shù)據(jù)，CUDA提供了原子操作函數(shù)——ATOM操作。
先亮出使用的例子：
[cpp] view plaincopy如果沒有加互斥機制，則同一個half?warp內(nèi)的線程將對i的操作混淆林亂。
用原子操作函數(shù)，可以很簡單的編寫自己的鎖，SKD中有給出的鎖結(jié)構(gòu)體如下：
[cpp] view plaincopy
8?CUDA軟件體系結(jié)構(gòu)
?
9?利用好現(xiàn)有的資源
如果連開方運算都需要自己去編寫程序?qū)崿F(xiàn)，那么我相信程序員這個職業(yè)將會縮水，沒有人愿意去干這種活。我想，程序員需要學(xué)會“偷懶”，現(xiàn)有的資源必須學(xué)會高效率的使用。當c++出現(xiàn)了STL庫，c++程序員的開發(fā)效率可以說倍增，而且程序穩(wěn)定性更高。
CUDA有提供給我們什么了嗎？給了，其實給了很多。
先介紹幾個庫：CUFFT、CUBLAS、CUDPP。
這里我先不詳細學(xué)習(xí)這些庫里到底有哪些函數(shù)，但是，大方向是需要了解的，不然找都不知道去哪兒找。CUFFT是傅里葉變換的庫，CUBLAS提供了基本的矩陣和向量運算，CUDPP提供了常用的并行排序、搜索等。
CUDA4.0以上，提供了一個類似STL的模板庫，初步窺探，只是一個類似vector的模板類型。有map嗎？map其實是一個散列表，可以用hashtable去實現(xiàn)這項機制。
SDK里面有很多例子，包括一些通用的基本操作，比如InitCUDA等，都可以固化成函數(shù)組件，供新程序的調(diào)用。
具體的一些可以固化的東西，我將在以后的學(xué)習(xí)中歸納總結(jié)，豐富自己的CUDA庫！
http://blog.csdn.net/huangfengxiao/article/details/8732789http://blog.csdn.net/huangfengxiao/article/details/8732790
總結(jié)
以上是生活随笔為你收集整理的CUDA: GPU高性能运算的全部內(nèi)容，希望文章能夠幫你解決所遇到的問題。
- 上一篇： 计算机用户组连接打印机,在组策略中使用脚
- 下一篇： java痴和堆_JAVA虚拟机理解 -