本小節筆記大綱：

• 1.Communication patterns
  • gather,scatter,stencil,transpose
• 2.GPU hardware & Programming Model
  • SMs,threads,blocks,ordering
  • Synchronization
  • Memory model: local, shared, global
    
  • Atomic Operation
• 3.Efficient GPU Programming
  • Access memory faster
    • coalescing global memory
    • use faster memory
  • Avoid thread divergence

一、Communication Patterns

1.Patterns

• Map

map很好理解，其實就是映射，也就是輸入和輸出一一對應，一個蘿卜一個坑

• Gather

Gather中文名為收集，是將若干個輸入數據經過計算后得到一個輸出值，如圖左示。很典型的應用就是比如說對于一個圖像，我們需要每一個像素值是其四周像素的平均值。

• Scatter
  scatter的特點是每個線程一次會向內存輸出多個值，也可能多個線程向一個內存輸出值。

• Stencil
  Stencil表示模板的意思，所以也就是計算的時候用模子來選擇輸入數據，看下圖就清楚了

• Transpose
  其實就是轉置啦~

具體應用實例如下：

在C語言中，加入我們定義了如上圖示的一個結構體，包含float和int兩種變量，然后我們又定義了一個該結構體的變量數組，一般來說其在內存中是像上面那樣排列的，強迫癥看起來是不是不舒服，而且這種排列方式比較浪費空間，所以通過轉置后形成下面的排列方式后既美觀又使運算加速了，豈不美哉？

2.練習題

• 第一個很簡單就是map，不仔細解釋了
• 第二個個表達式我之前腦袋一熱就選了C。。但是要注意，scatter的特點是每個線程一次會向內存輸出多個值，這顯然不符合該特點，而應該是Transpose。
• 第三個就是scatter了，原因如上
• 最后一個很容易選stencil，但是你要注意if條件語句的限制，所以應該是Gather。

3.總結神圖

二、GPU Hardware

1.問題導向

• 線程是如何有效地一致訪問內存
  • 子話題：如何利用數據重用
• 線程如何通過共享內存通信部分結果

2.硬件組成

如圖示，GPU由若干個SM(Stream Multiprocessor流多處理器)組成，而每個SM又包含若干個SP(教材上是Stream Processor流處理器，改視頻中是simple processor),anyway...開心就好，管他叫什么名字~

GPU的作用是負責分配線程塊在硬件SM上運行，所有SM都以并行獨立的方式運行。

下面做一下題目吧：

解析：

• 1正確.一個線程塊包含許多線程
• 2正確.一個SM可能會運行多個多個線程塊
• 3錯誤，因為一個線程塊無法在一個以上的SM上運行
• 4正確，在一個線程塊上所有線程有可能配合起來解決某個子問題
• 5錯誤，一個SM上可能有多個線程塊，但是根據定義，線程和不同的線程塊不應該存在協作關系。

3.程序員與GPU分工

另外需要注意的是程序員負責定義線程塊，而GPU則負責管理硬件，因此程序員不能指定線程塊的執行順序，也不能指定線程塊在某一特定的
SM上運行。

這樣設計的好處如下：

• 硬件可以運行的更加有效率
• 運行切換不需要等待，一旦一個線程塊運行完畢，SM可以自動的將另一個線程塊加載進來
• 最大的優勢：可擴展性，因為可以自動分配硬件資源，所以向下到單個SM，上到超級計算機的大量SM，均可以很好的適應。

有如上好處的同時，自然也就有局限性：

• 對于哪個塊在哪個SM上運行無法進行任何假設
• 無法獲得塊之間的明確的通信

4.GPU Memory Model

如圖示

• 每個線程都有它自己的本地內存(local memory)
• 線程塊有一個共享內存(shared memory)，塊中所有線程都可以訪問該內存中的數據
• GPU中的全局內存(global memory)是所有線程塊中的線程都能訪問的內存，也是CPU進行數據傳遞的地方。

訪問速度：

local memory > shared memory > global

例題：

解析：

s,t,u是本地內存中的變量，所以t=s最先運行，同理可以排除其他代碼運行順序。

注意：這只是為了說明訪問速度出的例題，實際情況中，編譯器可能會做出相應的調整來達到我們的目的

5.Sychronization

說道線程，很自然我們就需要考慮同步。GPU中的同步有如下幾種：

Barrier(屏障)

顧名思義，就是所有線程運行到這個點都需要停下來。

如圖示，紅色、藍色、綠色代表的線程先后到達barrier這個時間點后都停下來進行同步操作，完成之后線程的執行順序是不一定的，可能如圖示藍色線程先執行，綠色，紅色緊隨其后。

另外其實還有一種隱式的barrier，比如說先后啟動kernel A和kernel B，一般來說kernel B執行之前kernel A肯定是執行完畢了的。

說了這么多來做下題吧~233

題目：如下圖示，現在需要實現一個數組前移的操作，即后面一個往前面挪，共享數組大小是128，問為實現這個功能，需要設置幾次同步操作(或者說需要設置幾個barrier?)

解析：
最開始的時候沒想明白，寫了127,128，但是都不對。后來聽解釋才明白。前移操作可以分為三步：

• 為每個數組元素賦值，即

array[idx] = threadIdx.x; __syncthreads(); # 128個線程都執行完賦值語句后才能進行下一步

• 讀取后面一個元素的值，存在臨時變量里

int temp = array[idx+1]; __syncthreads();

• 將后一元素的值往前移

array[idx] = temp; __syncthreads();

6.Atomic Memory Operation

在cuda編程中經常會碰到這樣的情況，即大量的線程同時都需要對某一個內存地址進行讀寫操作，很自然這會發生沖突，如下圖示:

下面是發生沖突的具體的代碼示例：

#include "cuda_runtime.h" #include "device_launch_parameters.h" #include <stdio.h>#define NUM_THREADS 10000 #define ARRAY_SIZE 10 #define BLOCK_WIDTH 100void printDevice();__global__ void increment_naive(int *g) {int i = blockIdx.x * blockDim.x + threadIdx.x;i = i % ARRAY_SIZE;g[i] = g[i] + 1; }int main(int argc, char **argv) {printDevice();printf("\n");int h_array[ARRAY_SIZE];const int ARRAY_BYTES = ARRAY_SIZE * sizeof(int);int *d_array;// 分配內存cudaMalloc((void **) &d_array, ARRAY_BYTES);cudaMemset((void *) d_array, 0, ARRAY_BYTES);increment_naive<<<NUM_THREADS/BLOCK_WIDTH, BLOCK_WIDTH>>>(d_array);cudaMemcpy(h_array, d_array, ARRAY_BYTES, cudaMemcpyDeviceToHost);for(int i=0; i<ARRAY_SIZE; i++){printf("%d:%d\n",i,h_array[i]);}// 釋放內存cudaFree(d_array);getchar();//CUDA_SAFE_CALL(cudaGetDeviceCount(&deviceCount));return 0; }

運行結果：(每次運行的結果是不確定的)

這里就需要引入原子操作，只需要將讀寫函數進行如下修改

__global__ void increment_atomicNaive(int *g) {int i = blockIdx.x * blockDim.x + threadIdx.x;i = i % ARRAY_SIZE;atomicAdd(&g[i], 1); }

運行結果：

使用原子操作也是有一定限制的，如下：

• 只能使用一些特定的運算(如加、減、最小值、異或等運算，但是取模，求冪等運算則不行)和數據類型(一般是整型int)
• 每個線程塊里的不同線程以及線程塊本身將以不定的順序運行，我們在內存上用原子進行的運算順序也是不定的。
  例如下面的計算表達式的記過會不一樣：
  \(a+b+c 和 a+(b+c),其中a=1,b=10^,c=-10^{99}\)
• 雖然順序不確定，但是要知道的是GPU還是會強制每個線程輪流訪問內存，這把不同線程對內存的訪問串行化

提高CUDA編程效率策略

• 高運算密度(high arithmetic intensity)
  \(\frac{math}{memory}\)

前面提到了很多優化策略是集中在memory上的，把數據盡可能放到更快地內存上去，其中內存速度是
local > share > global

• 避免線程發散(avoid thread divergence)

如圖是線程發散的主要場景，即if else語句，上圖右邊非常生動的展現了線程發散的情形，可以看到各個線程在碰到if條件句后開始發散，最后聚合，但是最后各個線程之間的編號還是保持原來的不變的，這就是線程發散。

下面舉一個更加極端的例子，就是循環語句，如下圖示：

可以看到有藍、紅、綠、紫四個線程同時運行，藍線程只循環了一次，其他線程循環次數都多于藍線程，當藍線程退出循環后就不得不一直等著其他線程，上圖左下角的示意圖可以很直觀的看到這大大降低了運行效率，這也是為什么我們需要避免線程發散。

Summary

MARSGGBO原創





2017-11-1

總結

以上是生活随笔為你收集整理的Udacity并行计算课程笔记-The GPU Hardware and Parallel Communication Patterns的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。