CUDA: GPU高性能运算
CUDA: GPU高性能運算
2013-10-11 22:23 5650人閱讀 評論(0) 收藏 舉報 分類： CUDA（106）目錄(?)[+]
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，那多不過癮。
[cpp] view plaincopyI?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
5?GPU也不允許偏心
并行的事情多了，我們作為GPU的指令分配者，不能偏心了——給甲做的事情多，而乙沒事做，個么甲肯定不爽的來。所以，在GPU中，叫做線程網絡的分配。首先還是來看下GPU的線程網絡吧，圖2：
圖2?線程網絡
我們將具體點的，在主機函數中如果我們分配的是這樣的一個東西：
dim3?blocks(32, 32);
dim3?threads(16, 16);
dim3是神馬？dim3是一個內置的結構體，和linux下定義的線程結構體是個類似的意義的東西，dim3結構變量有x，y，z，表示3維的維度。不理解沒關系，慢慢看。
kernelfun<<<blocks,?threads>>>();
我們調用kernelfun這個內核函數，將blocks和threads傳到<<<,>>>里去，這句話可牛逼大了——相當于發號施令，命令那些線程去干活。這里使用了32*32?*?16*16個線程來干活。你看明白了嗎？blocks表示用了二維的32*32個block組，而每個block中又用了16*16的二維的thread組。好吧，我們這個施令動用了262144個線程！我們先不管GPU內部是如何調度這些線程的，反正我們這一句話就是用了這么多線程。
那我們的內核函數kernelfun()如何知道自己執行的是哪個線程？這就是線程網絡的特點啦，為什么叫網絡，是有講究的，網絡就可以定格到網點：
比如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
我習慣的用這樣的模式去分配，比較通用：
dim3?dimGrid();
dim3?dimBlock();
kerneladd<<<dimGrid,?dimBlock>>>();
這可是萬金油啊，你需要做的事情是填充dimGrid和dimBlock的結構體構造函數變量，比如，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）的網格編號。
我們這么理解吧，現在又一群人，我們分成16*16個小組（block），排列好，比如第3行第4列就指的是（2,3）這個小組。
而dimBlock（16,16）表示每個小組有16*16個成員，如果你想點名第3行第4列這個小組的里面的第3行第4列那個同學，那么，你就是在（2,3）這個block中選擇了（2,3）這個線程。這樣應該有那么一點可以理解進去的意思了吧？不理解透徹么什么關系，這個東西本來就是cuda中最讓我糾結的事情。我們且不管如何分配線程，能達到最優化，我們的目標是先讓GPU正確地跑起來，計算出結果即可，管他高效不高效，管他環保不環保。
嘮叨了這么多，下面我們用一個最能說明問題的例子來進一步理解線程網絡分配機制來了解線程網絡的使用。
一維網絡線程
eg：int?arr[1000]，對每個數組元素進行加1操作。
idea：我們最直接的想法，是調度1000個線程去干這件事情。
first?pro：我想用一個小組的1000個人員去干活。這里會存在這樣一個問題——一個小組是不是有這么多人員呢？是的，這個事情你必須了解，連自己組內多少人都不知道，你也不配作指揮官呀。對的，這個參數叫做maxThreadsPerBlock，如何取得呢？
好吧，cuda定義了一個結構體cudaDeviceProp，里面存入了一系列的結構體變量作為GPU的參數，出了maxThreadsPerBlock，還有很多信息哦，我們用到了再說。
maxThreadsPerBlock這個參數值是隨著GPU級別有遞增的，早起的顯卡可能512個線程，我的GT520可以跑1024個線程，辦公室的GTX650ti2G可以跑1536個，無可非議，當然多多益善。一開始，我在想，是不是程序將每個block開的線程開滿是最好的呢？這個問題留在以后在說，一口吃不成胖子啦。
好吧，我們的數組元素1000個，是可以在一個block中干完的。
內核函數：
[cpp] view plaincopy
呀，原來這么簡單，個么CUDA也忒簡單了哇！這中想法是好的，給自己提高信心，但是這種想法多了是不好的，因為后面的問題多了去了。
盆友說，1000個元素，還不如CPU來的快，對的，很多情況下，數據量并行度不是特別大的情況下，可能CPU來的更快一些，比較設備與主機之間互相調度操作，是會有額外開銷的。有人就問了，一個10000個元素的數組是不是上面提供的idea就解決不了啦？對，一個block人都沒怎么多，如何完成！這個情況下有兩條路可以選擇——
第一，我就用一個組的1000人來干活話，每個人讓他干10個元素好了。
這個解決方案，我們需要修改的是內核函數：
[cpp] view plaincopy第二，我多用幾個組來干這件事情，比如我用10個組，每個組用1000人。
這個解決方案就稍微復雜了一點，注意只是一點點哦~因為，組內部怎么干活和最原始的做法是一樣的，不同之處是，我們調遣了10個組去干這件事情。
首先我們來修改我們的主機函數：
[cpp] view plaincopy盆友要問了，10個組每個組1000人，你怎么點兵呢？很簡單啊，第1組第3個線程出列，第9組第9個線程出列。每個人用組號和組內的編號定了位置。在線程網絡中，blockId.x和threadId.x就是對應的組號和組內編號啦，我必須要這里開始形象點表示這個對應關系，如果這個對應關系是這樣子的[blockId.x，threadId.x]，那么我們的數組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]
是不是很有規律呢？對的，用blockId.x和threadId.x可以很好的知道哪個線程干哪個元素，這個元素的下表就是threadId.x?+?1000*blockId.x。
這里我想說的是，如果我們哪天糊涂了，畫一畫這個對應關系的表，也許，就更加清楚的知道我們分配的線程對應的處理那些東西啦。?
一維線程網絡，就先學這么多了。
二維網絡線程
eg2：int?arr[32][16]二維的數組自增1。
第一個念頭，開個32*16個線程好了哇，萬事大吉！好吧。但是，朕現在想用二維線程網絡來解決，因為朕覺得一個二維的網絡去映射一個二維的數組，朕看的更加明了，看不清楚自己的士兵，如何帶兵打仗！
我還是畫個映射關系：
一個block中，現在是一個二維的thread網絡，如果我用了16*16個線程。
(0,0)，(0,1)，……(0,15)
(1,0)，(1,1)，……(1,15)
……
(15,0)，(15,1)，……(15,15)
呀，現在一個組內的人稱呼變了嘛，一維網絡中，你走到一個小組里，叫3號出列，就出來一個，你現在只是叫3號，沒人會出來！這個場景是這樣的，現在你班上有兩個人同名的人，你只叫名，他們不知道叫誰，你必須叫完整點，把他們的姓也叫出來。所以，二維網絡中的(0,3)就是原來一維網絡中的3，二維中的(i,j)就是一維中的(j+i*16)。不管怎么樣，一個block里面能處理的線程數量總和還是不變的。
一個grid中，block也可以是二維的，一個block中已經用了16*16的thread了，那我們一共就32*16個元素，我們用2個block就行了。
先給出一個代碼清單吧，程序員都喜歡看代碼，這段代碼是我抄襲的。第一次這么完整的放上代碼，因為我覺得這個代碼可以讓我說明我想說的幾個問題：
第一，二維數組和二維指針的聯系。
第二，二維線程網絡。
第三，cuda的一些內存操作，和返回值的判斷。
[cpp] view plaincopy簡要的來學習一下二維網絡這個知識點，?
dim3?dimBlock(16, 16);?
//定義block內的thread二維網絡為16*16
dim3?dimGrid((COLS + dimBlock.x - 1) / (dimBlock.x), ?(ROWS + dimBlock.y - 1) / (dimBlock.y));?
//定義grid內的block二維網絡為1*2
unsigned?int?row?=?blockDim.y * blockIdx.y ?+ ?threadIdx.y;?
//二維數組中的行號
unsigned?int?col?=?blockDim.x * blockIdx.x ?+ ?threadIdx.x;?
//二維線程中的列號
三維網絡線程
dim3定義了三維的結構，但是，貌似二維之內就能處理很多事情啦，所以，我放棄學習三維。網上看到的不支持三維網絡是什么意思呢？先放一放。
給自己充充電
同一塊顯卡，不管你是二維和三維或一維，其計算能力是固定的。比如一個block能處理1024個線程，那么，一維和二維線程網絡是不是處理的線程數一樣呢？
?
回答此問題，先給出網絡配置的參數形式——<<<Dg,Db,Ns,S>>>，各個參數含義如下：
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：這個是可選參數，設定最多能動態分配的共享內存大小，比如16k，單不需要是，這個值可以省略或寫0。
S：也是可選參數，表示流號，默認為0。流這個概念我們這里不說。
接著，我想解決幾個你肯定想問的兩個問題，因為我看很多人想我這樣的問這個問題：
1?block內的thread我們是都飽和使用嗎？
答：不要，一般來說，我們開128或256個線程，二維的話就是16*16。
2?grid內一般用幾個block呢？
答：牛人告訴我，一般來說是你的流處理器的4倍以上，這樣效率最高。
回答這兩個問題的解釋，我想抄襲牛人的一段解釋，解釋的好的東西就要推廣呀：
GPU的計算核心是以一定數量的Streaming?Processor(SP)組成的處理器陣列，NV稱之為Texture?Processing?Clusters(TPC)，每個TPC中又包含一定數量的Streaming?Multi-Processor(SM)，每個SM包含8個SP。SP的主要結構為一個ALU（邏輯運算單元），一個FPU（浮點運算單元）以及一個Register?File(寄存器堆)。SM內包含有一個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模型中最基本的運行單元，執行最基本的程序指令。Block是一組協作Thread，Block內部允許共享存儲，每個Block最多包含512個Thread。Grid是一組Block，共享全局內存。Kernel是在GPU上執行的核心程序，每一個Grid對應一個Kernel任務。 在程序運行的時候，實際上每32個Thread組成一個Warp，每個?warp?塊都包含連續的線程，遞增線程?ID?。Warp是MP的基本調度單位，每次運行的時候，由于MP數量不同，所以一個Block內的所有Thread不一定全部同時運行，但是每個Warp內的所有Thread一定同時運行。因此，我們在定義Block?Size的時候應使其為Warp?Size的整數倍，也就是Block?Size應為32的整數倍。理論上Thread越多，就越能彌補單個Thread讀取數據的latency?，但是當Thread越多，每個Thread可用的寄存器也就越少，嚴重的時候甚至能造成Kernel無法啟動。因此每個Block最少應包含64個Thread，一般選擇128或者256，具體視MP數目而定。一個MP最多可以同時運行768個Thread，但每個MP最多包含8個Block，因此要保持100%利用率，Block數目與其Size有如下幾種設定方式： ??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?
這些電很重要啊，必須要充！不然，我就很難理解為什么網絡線程如何分配的。
6?規約思想和同步概念
擴大點說，并行計算是有一種基本思想的，這個算法能解決很多很常規的問題，而且很實用，比如說累加和累積等——規約思想。對于基礎的、重要的，我想有必要系統的學習。
我覺得有必要重新復制下之前寫的這篇介紹：
http://www.cnblogs.com/viviman/archive/2012/11/21/2780286.html
并行程序的開發有其不同于單核程序的特殊性，算法是重中之重。根據不同業務設計出不同的并行算法，直接影響到程序的效率。因此，如何設計并行程序的算法，似乎成為并編程的最大難點。觀其算法，包括cuda?sdk的例子和網上的牛人，給出的一些例子，以矩陣和矢量處理為主，深入點的包括fft和julia等數學公式，再高級一點的算是圖形處理方面的例子。學習這些算法的思想，免不了有自己的一點點總結。之前學習過omp編程，結合現在的cuda，我覺得要理解并行編程，首先理解劃分和規約這兩個概念。也許你的算法學的更加扎實。劃分是《算法》里面的一個重要思想，將一個大的問題或任務，分解成小問題小任務，各個擊破，最后歸并結果；規約是《cuda**》書上介紹的一個入門的重要思想，規約算法(reduction)用來求連加、連乘、最值等，應用廣泛。每次循環參加運算的線程減少一半。不管算法的思想如何花樣，萬變不離其中的一點--將一個大的任務分解成小的任務集合，分解原則是粒度合適盡量小、數據相關性盡量小。如此而已。因為，我們用GPU是為了加速，要加速必須提高執行任務的并行度！明白這個道理，那么我們將絞盡腦汁地去想方設法分析自己手上的任務，分解、分解、分解！這里拿規約來說事情，因為，規約這個東西，似乎可以拿來單做9*9乘法表來熟悉，熟悉了基礎的口訣，那么99*99的難題也會迎刃而解。ex：矢量加法，要實現N=64*256長度的矢量的累加和。假設a+b計算一次耗時t。
cpu計算：顯然單核的話需要64*256*t。我們容忍不了。
gpu計算：最初的設想，我們如果有個gpu能同時跑N/2個線程，我們這N/2個線程同時跑，那么不是只需要t時間就能將N個數相加編程N/2個數相加了嗎？對的。這一輪我們用了t時間；接著的想法，我們不斷的遞歸這個過程，能發現嗎？第二輪，我們用N/2/2個線程同時跑，剩下N/2/2個數相加，這一輪我們同樣用了t時間；一直這樣想下去吧，最后一輪，我們用了1個線程跑，剩下1個數啦，這就是我們的結果！每一輪時間都為t，那么理想情況，我們用了多少輪這樣的計算呢？計算次數=log(N)=6*8=48，對的，只用了48輪，也就是說，我們花了48*t的時間！
規約就是這樣，很簡單，很好用，我們且不管程序后期的優化，單從這個算法分析上來說，從時間復雜度N降到了logN，這在常規算法上，要提高成這樣的效率，是不得了的，這是指數級別的效率提高！所以，你會發現，GPU有CPU無法取代的得天獨厚的優勢——處理單元真心多啊！
規約求和的核函數代碼如下：
[cpp] view plaincopy這個例子還讓我學到另一個東西——同步！我先不說同步是什么，你聽我說個故事：我們調遣了10個小組從南京去日本打仗，我們的約定是，10個組可以自己行動，所有組在第三天在上海機場會合，然后一起去日本。這件事情肯定是需要處理的，不能第1組到了上海就先去日本了，這些先到的組，唯一可以做的事情是——等待！這個先來后到的事情，需要統一管理的時候，必須同步一下，在上海這個地方，大家統一下步調，快的組等等慢的組，然后一起干接下去的旅程。
是不是很好理解，這就是同步在生活中的例子，應該這樣說，計算機的所有機制和算法很多都是源于生活！結合起來，理解起來會簡單一點。
在CUDA中，我們的同步機制用處大嗎？又是如何用的呢？我告訴你，一個正常規模的工程中，一般來說數據都會有先來后到的關系，這一個計算結果可能是提供給另一個線程用的，這種依賴關系存在，會造成同步的應用。
__synctheads()這句話的作用是，這個block中的所有線程都執行到此的時候，都聽下來，等所有都執行到這個地方的時候，再往下執行。
7?撬開編程的鎖
鎖是數據相關性里面肯定要用到的東西，很好，生活中也一樣，沒鎖，家里不安全；GPU中沒鎖，數據會被“盜”。
對于存在競爭的數據，CUDA提供了原子操作函數——ATOM操作。
先亮出使用的例子：
[cpp] view plaincopy如果沒有加互斥機制，則同一個half?warp內的線程將對i的操作混淆林亂。
用原子操作函數，可以很簡單的編寫自己的鎖，SKD中有給出的鎖結構體如下：
[cpp] view plaincopy
8?CUDA軟件體系結構
?
9?利用好現有的資源
如果連開方運算都需要自己去編寫程序實現，那么我相信程序員這個職業將會縮水，沒有人愿意去干這種活。我想，程序員需要學會“偷懶”，現有的資源必須學會高效率的使用。當c++出現了STL庫，c++程序員的開發效率可以說倍增，而且程序穩定性更高。
CUDA有提供給我們什么了嗎？給了，其實給了很多。
先介紹幾個庫：CUFFT、CUBLAS、CUDPP。
這里我先不詳細學習這些庫里到底有哪些函數，但是，大方向是需要了解的，不然找都不知道去哪兒找。CUFFT是傅里葉變換的庫，CUBLAS提供了基本的矩陣和向量運算，CUDPP提供了常用的并行排序、搜索等。
CUDA4.0以上，提供了一個類似STL的模板庫，初步窺探，只是一個類似vector的模板類型。有map嗎？map其實是一個散列表，可以用hashtable去實現這項機制。
SDK里面有很多例子，包括一些通用的基本操作，比如InitCUDA等，都可以固化成函數組件，供新程序的調用。
具體的一些可以固化的東西，我將在以后的學習中歸納總結，豐富自己的CUDA庫！
http://blog.csdn.net/huangfengxiao/article/details/8732789http://blog.csdn.net/huangfengxiao/article/details/8732790
總結
以上是生活随笔為你收集整理的CUDA: GPU高性能运算的全部內容，希望文章能夠幫你解決所遇到的問題。
- 上一篇： 计算机用户组连接打印机,在组策略中使用脚
- 下一篇： java痴和堆_JAVA虚拟机理解 -