前言

《并行編程》系列是學(xué)習(xí)《Intro to Parallel Programming》過程中所做的筆記記錄以及個人一些所思所想希柿。

并行通信

并行計算需要解決的一個問題就是泞歉，如何解決線程之間的協(xié)同工作（Working together）問題。而協(xié)同工作的關(guān)鍵則是通信（Communication）。

CUDA 的通信發(fā)生在內(nèi)存當(dāng)中甲雅，例如履怯，多個線程可能需要從同一個內(nèi)存地址中讀取數(shù)據(jù)艇潭；也可能出現(xiàn)多個線程需要同時向同一個地址寫入數(shù)據(jù)拼窥；可能出現(xiàn)多個線程之間需要交換數(shù)據(jù)。

并行通信模式（Parallel Communication Patterns）

并行通信存在多種模式蹋凝，通信模式反映了線程所執(zhí)行的任務(wù)與內(nèi)存之間的映射關(guān)系鲁纠。這里將分別介紹五種通信模式：

1. Map
2. Gather
3. Scatter
4. Stencil
5. Transpose

Map

Map: Tasks read from and write to specific data elements.

Map 模式下，每個線程將從內(nèi)存的特定地址中讀取數(shù)據(jù)進(jìn)行處理仙粱，然后再寫入特定的地址中，它的輸入與輸出具有嚴(yán)格的一對一的關(guān)系彻舰。

Map 模式在 GPU 中非常高效伐割，在 CUDA 中也能很容易通過有效的方式表達(dá)。

但是 Map 比較不靈活刃唤，能處理的問題有限隔心。

Gather

現(xiàn)在假設(shè)需要求取3個數(shù)據(jù)的平均值，那么在 Gather 模式下尚胞，每個線程將從內(nèi)存中的三個位置讀取數(shù)據(jù)硬霍，然后將這三個數(shù)取平均，寫入指定的內(nèi)存中笼裳。

這一模式可用于涉及到濾波器的一系列操作唯卖。

Scatter

Scatter: Tasks compute where to write output.

與 Gather 模式下粱玲，多個輸入一個輸出相反，Scatter 模式是一個輸入多個輸出拜轨。

另外在同時寫入多個輸出的時候?qū)⒊霈F(xiàn)沖突的問題抽减，這將在后續(xù)進(jìn)行討論。

Stencil

Stencil: Tasks read input from a fixed neighborhood in an array.

常用的模板有：

1. 2D von Neumann

   
   

2. 2D Moore

   
   

3. 3D von Neumann

   
   

看到這里橄碾，可能有人會對 Stencil 和 Gather 產(chǎn)生疑惑卵沉。咋看之下，兩者確實非常相似法牲，但是 Stencil 模式中史汗，要求每個線程都嚴(yán)格執(zhí)行相同的模板，但是 Gather 模式卻沒有這個限制拒垃，因此停撞，比如說，在 Gather 模式中就可以按線程索引的奇偶不同恶复，給線程分配不同的操作任務(wù)怜森。

Transpose

Transpose: Tasks re-order data elements in memory.

對于一張圖像，其數(shù)據(jù)在內(nèi)存中的存儲的方式如下：

但是在某些情況下谤牡，可能需要將圖像轉(zhuǎn)置副硅。

通常在涉及到數(shù)組運算、矩陣運算和圖像操作的時候會需要使用到 Transpose翅萤，但是 Transpose 也適用于其它數(shù)據(jù)結(jié)構(gòu)恐疲。

比如定義了一個結(jié)構(gòu)體 foo ，然后創(chuàng)建一個該結(jié)構(gòu)的結(jié)構(gòu)數(shù)組（AoS）套么，如果想將該結(jié)構(gòu)數(shù)組變換成數(shù)組結(jié)構(gòu)（SoA）培己，也可以通過 Transpose 實現(xiàn)。

總結(jié)

上圖總結(jié)了并行計算的七種計算模式胚泌，除了之前介紹的五種模式以外省咨，還有兩種更加基礎(chǔ)的模式將在接下來進(jìn)行介紹。

GPU

程序員眼中的 GPU

程序員在并行編程中所要做的就是玷室，創(chuàng)建內(nèi)核（C/C++函數(shù)）用來處理具體的任務(wù)零蓉。內(nèi)核由許多線程（完整執(zhí)行一段處理程序的通路）組成，圖中的線程都采用曲線繪制穷缤，其原因是敌蜂，每個線程的具體通路可能不相同（即每個線程所執(zhí)行的運算不相同）。

多個線程將組成線程塊津肛，一個線程塊內(nèi)的多個線程負(fù)責(zé)協(xié)同處理一項任務(wù)或者子任務(wù)章喉。

上圖中，程序首先啟動了一個內(nèi)核 foo，等到其中所有的線程都運行完了之后秸脱，結(jié)束內(nèi)核落包。然后又啟動了內(nèi)核 bar，可以注意到撞反，一個內(nèi)核中所具有的線程塊妥色，以及每個線程塊中的線程數(shù)是可以自己配置的參數(shù)。

線程塊與 GPU

GPU 中包含有許多的流處理器（Streaming Multiprocessor, SM）遏片，不同的 GPU 包含有不同數(shù)量的流處理器嘹害，并且流處理器數(shù)量也是衡量 GPU 性能的一項重要指標(biāo)。

一個流處理器中包含有多個簡單的處理器和內(nèi)存吮便。

當(dāng)你的程序創(chuàng)建了內(nèi)核之后笔呀，GPU 將為內(nèi)核中的線程塊分配流處理器，每個線程塊被分配給一個流處理器髓需，然后這些流處理器以并行的方式進(jìn)行運行许师。

注意：一個流處理器上允許運行多個線程塊，但是一個線程塊只允許被分配給一個流處理器運行僚匆。

CUDA 特征

CUDA 不具備的特征

CUDA 對于內(nèi)核中的線程塊要何時運行微渠、該如何分配流處理以及有多少線程塊需要同時運行等細(xì)節(jié)沒有進(jìn)行任何的控制，這些分配問題都交給 GPU 進(jìn)行控制咧擂。這么做的好處有：

• 硬件將可以更加高效地執(zhí)行計算
• 當(dāng)一個線程塊執(zhí)行完成之后逞盆，當(dāng)前的流處理器馬上又可以任意執(zhí)行下一個線程塊
• 更高的擴展性。因為流處理器的分配交由硬件控制松申，所以程序可以很好地在具有不同流處理器數(shù)量的設(shè)備上進(jìn)行移植云芦。

但是 CUDA 的這種做法也將導(dǎo)致一些后果：

• 對于某一線程塊將在哪個流處理器上運行無法做出任何預(yù)測
• 線程塊之間沒有通信交流。如果線程塊 x 的輸入依賴于線程塊 y 的輸出贸桶，而 y 已經(jīng)完成執(zhí)行并且退出舅逸，這將導(dǎo)致 x 的計算出現(xiàn)問題。這種現(xiàn)象稱為“dead lock”
• 線程塊中的線程不能永遠(yuǎn)執(zhí)行（比如皇筛，死循環(huán)）琉历，因為它需要在執(zhí)行完成之后釋放流處理器資源，以便于其它線程塊可以使用

CUDA 具備的特征

CUDA 在程序運行的時候水醋，能夠保證兩點：

• 同一個線程塊上的所有線程將同時在同一個流處理器中運行旗笔。
• 下一個內(nèi)核中的線程塊必須等待當(dāng)前內(nèi)核中的所有塊運行完成之后，才能運行离例。
  • 比如說换团，程序依次定義了兩個內(nèi)核 foo 和 bar 悉稠，bar 中的線程塊只有等到 foo 中的所有線程塊都運行完之后才能開始運行宫蛆。

GPU 內(nèi)存模型

每個線程都擁有一個局部內(nèi)存（Local memory），這就好像局部變量一樣，只有對應(yīng)的線程才能訪問耀盗。

然后想虎，線程塊也有一塊對應(yīng)的共享內(nèi)存（Shared memory）。共享內(nèi)存只能被對應(yīng)線程塊內(nèi)的線程進(jìn)行訪問叛拷。

另外還有具有全局內(nèi)存（Global memory）舌厨。不僅內(nèi)核中的所有線程可以訪問它，不同內(nèi)核也可以進(jìn)行訪問忿薇。

前邊介紹的局部內(nèi)存裙椭、共享內(nèi)存和全局內(nèi)存都是屬于 GPU 內(nèi)部的內(nèi)存。上圖展示了署浩，CPU 的線程啟動了 GPU 揉燃，然后將主機內(nèi)存（Host memory）中的數(shù)據(jù)拷貝到 GPU 的全局內(nèi)存中，以便于 GPU 內(nèi)核線程可以訪問這些數(shù)據(jù)筋栋。另外 GPU 內(nèi)核線程也可以直接訪問主機內(nèi)存炊汤，這一點將在后邊介紹。

同步

通過共享內(nèi)存和全局內(nèi)存弊攘，線程之間可以互相訪問彼此的計算結(jié)果抢腐，這也意味著線程間可以進(jìn)行協(xié)同計算。但是這樣也存在著風(fēng)險襟交， 如果一個線程在另一個線程寫入數(shù)據(jù)之前就讀取了數(shù)據(jù)怎么辦迈倍？因此線程之間需要同步的機制，來避免這種情形出現(xiàn)婿着。

事實上授瘦，同步問題是并行計算的一個最基本的問題。而解決同步問題的一個最簡單方法則是屏障（Barrier）竟宋。

Barrier: Point in the program where threads stop and wait. When all threads have reached the barrier, they can proceed.

屏障語句是 __syncthreads() 提完。

編程模型

現(xiàn)在，可以重新構(gòu)建一下編程模型丘侠。我們擁有線程和線程塊徒欣，并且在線程塊內(nèi)，可以創(chuàng)建屏障用于同步線程蜗字。事實上打肝，如果一個程序中創(chuàng)建了多個內(nèi)核，內(nèi)核之間默認(rèn)具有隱性的屏障挪捕，這使得不會出現(xiàn)多個內(nèi)核同時運行的情況粗梭。

然后再將之前介紹的內(nèi)存模型添加進(jìn)來，便得到了 CUDA 级零。

因此断医，CUDA 的核心就是層級計算結(jié)構(gòu)。從線程到線程塊再到內(nèi)核，對應(yīng)著內(nèi)存空間中的局部內(nèi)存鉴嗤、共享內(nèi)存和全局內(nèi)存斩启。

編寫高效的 CUDA 程序

這里將首先從頂層的策略上介紹如何編寫高效的 CUDA 程序。

首先需要知道的是 GPU 具有非常驚人的計算能力醉锅，一個高端的 GPU 可以實現(xiàn)每秒超過 3 萬億次的數(shù)學(xué)運算（3 TFLOPS/s）兔簇。但是如果一個 CUDA 程序的大多數(shù)時間都花費在了等待內(nèi)存的讀取或?qū)懭氩僮鞯脑挘@就相當(dāng)浪費計算能力硬耍。所以要編寫高效的 CUDA 程序的第一點是——最大化計算強度垄琐。

計算強度表達(dá)為每個線程計算操作時間除以每個線程在的訪存時間。所以要最大化計算強度经柴，就可以通過最大化分子和最小化分母來實現(xiàn)此虑。然而由于計算操作時間主要受具體算法的計算量限制，所以為了最大化計算強度主要從最小化訪存時間入手口锭。

最小化訪存時間

要最小化訪存時間的一種方式就是朦前，將訪問頻率更高的數(shù)據(jù)移動到訪問速度更快的內(nèi)存中。

在之前的介紹當(dāng)中已經(jīng)了解了 GPU 線程可以訪問四種類型的內(nèi)存鹃操，其中最快就是局部內(nèi)存韭寸。

局部內(nèi)存

局部變量的定義是最簡單的。

對于上圖的內(nèi)核代碼荆隘，變量 f 與參數(shù) in 都將存儲于局部內(nèi)存中恩伺。

共享內(nèi)存

要定義存儲于共享內(nèi)存中的變量，需要在變量定義語句前加一個 __shared__ 關(guān)鍵字進(jìn)行修飾椰拒。定義于共享內(nèi)存中的變量可以被同一個線程塊中的所有線程所訪問晶渠，其生存時間為線程塊的生存時間。

全局內(nèi)存

全局的內(nèi)存訪問要稍微麻煩些燃观，但是可以通過指針的機制來實現(xiàn)褒脯。

這里傳入內(nèi)核的參數(shù)被定義成一個指針，而這個指針恰恰指向的是全局內(nèi)存區(qū)域缆毁。

然后在 CPU 的代碼部分番川，首先創(chuàng)建了一個長度為 128 的浮點數(shù)數(shù)組 h_arr，它將存儲于主機內(nèi)存中（這里通過前綴 h_ 表明當(dāng)前變量運行于 HOST 中）脊框，然后定義了一個指向 GPU 全局內(nèi)存的指針 d_arr颁督，并通過 cudaMalloc 函數(shù)為 d_arr 分配全局存儲區(qū)域。

最小化訪存時間的另一個方法是使用合并全局內(nèi)存訪問（Coalesce global memory accesses）浇雹。

單一線程在訪問內(nèi)存時具有一個特性沉御，就是即使該線程只需要使用到內(nèi)存中的一小部分，但是程序也會從內(nèi)存中讀取一段連續(xù)的內(nèi)存塊昭灵。因此吠裆，如果此時恰好有其它線程也在使用該內(nèi)存塊中的數(shù)據(jù)聂儒，內(nèi)存塊就得到復(fù)用，從而節(jié)省再次讀取內(nèi)存的時間硫痰。

所以如果多個線程同時讀取或者寫入連續(xù)的全局內(nèi)存位置，此時 GPU 的效率的是最高的窜护，而這種訪問模式被稱為合并（Coalesced）效斑。

但是當(dāng)多個線程所訪問的全局內(nèi)存位置不連續(xù)或者甚至隨機的時候，此時 GPU 便無法繼續(xù)保持高效柱徙，因為很可能需要分別讀取全局內(nèi)存中的多個塊缓屠，這樣就增加了訪存時間。

相關(guān)性問題（Related problem）

Related problem: lots of threads reading and writing same memory locations

當(dāng)多個線程同時參與到對同一塊內(nèi)存地址的讀寫操作時护侮，將引發(fā)沖突從而導(dǎo)致錯誤的計算結(jié)果敌完，這便是相關(guān)性問題。

解決該相關(guān)性問題的一個方法是使用原子內(nèi)存操作（Atomic memory operations）羊初。

原子內(nèi)存操作

CUDA 提供了若干個原子內(nèi)存操作函數(shù)滨溉，通過這些函數(shù)可以以原子操作的方式訪問內(nèi)存，也就是某一時刻內(nèi)存中的特定地址只能被單一線程所讀寫长赞，從而避免了相關(guān)性問題晦攒。

常見的原子內(nèi)存操作：

• atomicAdd()，原子相加
• atomicMin()得哆，原子最小值
• atomicXOR()脯颜，原子異或
• atomicCAS()，比較并且交換（Compare-and-Swap）

說明：這些原子內(nèi)存操作函數(shù)的實現(xiàn)借助了硬件來實現(xiàn)原子操作贩据，這里將不進(jìn)行介紹栋操。

但是這些原子操作也存在一些局限性。

• 只支持某些特定的操作（比如饱亮，支持加矾芙、減、最小值和異或等近上，不支持求余蠕啄、求冪等操作）和數(shù)據(jù)類型（主要支持整數(shù)）。
• 沒有順序限制戈锻。盡管使用了原子操作歼跟，但是關(guān)于線程執(zhí)行順序的問題依然沒有定義。
  • 由于浮點數(shù)精度問題格遭，這將導(dǎo)致浮點數(shù)運算出現(xiàn)非關(guān)聯(lián)現(xiàn)象（Non-associative）哈街。具體來說就是可能出現(xiàn) (a + b) + c != a + (b + c)，比如拒迅，當(dāng) a = 1, b = 10^99, c= 10^-99 時骚秦。
• 串行化線程內(nèi)存訪問她倘。原子操作的實現(xiàn)并沒有使用什么神奇的魔法，它僅僅只是串行化了線程對同一個內(nèi)存地址的訪問作箍，所以這將減慢整體的計算速度硬梁。

線程發(fā)散

前邊已經(jīng)介紹過了，要使得 CUDA 程序高效的一個關(guān)鍵點是——最大化計算強度胞得。然后另外一個關(guān)鍵點是——避免線程發(fā)散（Thread divergence）荧止。

線程發(fā)散指的是，比如說當(dāng)內(nèi)核代碼中出現(xiàn)條件語句時阶剑，線程運行到條件語句處跃巡，可能有些線程符合條件，而有些線程不符合條件牧愁，此時它們就會發(fā)散開素邪，形成兩條路徑，然后在條件語句塊結(jié)束之后再次聚合到同一條路徑上猪半。

不僅僅只有條件語句才會導(dǎo)致線程發(fā)散兔朦，循環(huán)語句也可能導(dǎo)致。

舉個不太恰當(dāng)?shù)睦幽ト罚谶@個內(nèi)核代碼中有一個循環(huán)烘绽，循環(huán)的次數(shù)是當(dāng)前線程的索引。

所以線程的執(zhí)行路徑如上圖俐填，如果以時間為橫軸繪制線程運行圖則如下圖安接。

由于硬件傾向于同時執(zhí)行完線程，所以當(dāng)線程索引小的線程完成循環(huán)之后英融，它還會繼續(xù)等待其它線程完成循環(huán)盏檐，直至所有線程都完成循環(huán)之后，這些線程才會繼續(xù)執(zhí)行循環(huán)塊之后的代碼驶悟。因此胡野，這里除了最后一個線程充分利用了時間進(jìn)行運算以外，其它線程均無法有效利用時間痕鳍。而這也就是為什么要避免線程發(fā)散的原因硫豆。

總結(jié)

本節(jié)內(nèi)容小結(jié)：

• 通信模式
  • gather, scatter, stencil, transpose
• GPU 硬件與編程模型
  • 流處理器，線程笼呆，線程塊
  • 線程同步
  • 內(nèi)存模型（局部熊响，共享，全局诗赌，主機）汗茄，原子操作
• 高效 GPU 編程
  • 減少訪存花銷（使用更快的內(nèi)存，合并全局內(nèi)存訪問）
  • 避免線程發(fā)散

課堂作業(yè)

本次的課堂作業(yè)是實現(xiàn)圖像模糊铭若，思路相對較簡單洪碳。唯一需要注意的是邊界情況的取值递览。因為當(dāng) filter 的中心位于圖像邊界的時候，它的周圍像素會出現(xiàn)超出圖像的現(xiàn)象瞳腌，這里需要進(jìn)行判斷绞铃。

課程作業(yè)完成代碼：
https://github.com/un-knight/cs344-parallel-programming