GPU硬件结构和程序具体参数设置
目錄
1、CUDA程序Thread的基本結構
3、單指令多線程架構（SIMT, Single Instruction Multi Thread）
4、硬件多線程（Hardward Multithreading）
5、GPU的顯存結構（Memory Hierarchy）
本文主要對GPU的硬件，以及根據硬件定量對參數進行設置，按照先了解硬件在進行參數設置的順序分別進行描述。
1、CUDA程序Thread的基本結構
在邏輯上，threads分為如下三個層次：
而在硬件上，threads僅有兩個層次：
2、流處理器（SM）
GPU由多個multithreaded Streaming Multiprocessos(SM)構成。
以Tesla K40c顯卡為例（Compute Capability 3.5）為例，其有15個SM。每個SM包括192個CUDA cores。
當要執行一個kernel grid時，該grid中的blocks會被分配給可用的SM。
如果一個block執行完畢，那么新的block會被放到空出來的SM上執行。
SM的設計目的時并發執行幾百個線程，因此使用了SIMT, Single Instruction Multi Thread?的架構。
在SM內部有兩種級別的并行：
3、單指令多線程架構（SIMT, Single Instruction Multi Thread）
在SM內部，threads以warp為單位被創建/管理/調度和執行，每個warp包括32個threads。當將一個或多個blocks分配給SM時，它會首先將其分成多個warp。（每個warp所包含的thread都是按threadID有序遞增的），然后使用warp scheduler來調度執行每個warp。一個warp內的32個threads在同一時間執行同一條指令，所以當32個thread的執行路徑完全一致時效率最高。如果有data-dependent的分支，那么warp會分別執行每一個分支路徑，不在當前分支的threads會被停用。
4、硬件多線程（Hardward Multithreading）
每個warp的執行上下文execution context（program counters, registers, etc.）在整個生命周期里都是被保存在片上存儲（on-chip memory）上的，但是片上存儲大小是比較小的。片上存儲可以類似于CPU/單片機的內部寄存器之類的硬件，因為片上內存有限，所以現在計算機中都是包含有片外存儲的。因此從一個execution context切換到另一個execution context是無消耗的，在每個instruction issue time里，一個warp scheduler都會選擇一個warp，該warp中的threads需要做好準備執行下個instruction，然后給向這個warp里的threads發出指令。
具體而言，每個Multiprocessor都有
- 一組32-bit registers（按照warp數來分配）
- 一個【parallel data cache/shared memory】（按照thread blocks數來分配）
這兩個條件就決定了一個SM上能同時【并發的】存在多少個warps和blocks。（同時也有最大值限制）。如果一個block需要的registers/shared mems都無法滿足，那么kernel就會失敗。更細節一些，即在每個instruction issue time，一個warp scheduler都會選擇一個準備好的warp發出指令。等待warp準備好的這段時間（number of clock cycles）就是【latency】。要達到完全的利用率，就需要所有的warp scheduler在latency這段時間的每個clock cycles都可以發出指令給其他warp，即掩蓋掉latency。因此，一個SM內越多的warp通常就會帶來越高的利用率，性能越高。
5、GPU的顯存結構（Memory Hierarchy）
在了解GPU的內存層次之前，我們先了解下如下術語：
- Cache Line：每次讀或寫內存時，即使只操作一個值，也是會把一小塊內存讀取到Cache里的。這一小塊被讀取到Cache的內存就叫【Cache Line】，其大小稱之為【Cache Line Size】
- Memory Transaction：a transaction is the movement of a unit of data between two regions of memory。例如，從Mem到L2 Cache的一次拷貝就是一次Mem Transaction。
- Register Spilling：某些應該放到register的變量，由于register不夠大，而放到了mem中（GPU中是放在Local Mem）
- Natrually Aligned：any item is aligned to at least a multiple of its own size。例如4Byte的對象的地址必須能整除4；8Byte的對象的地址必須能整除8
GPU的顯存層次如下：
全局顯存（Global Memory）
Global Memory就是編寫CUDA程序時最常使用的顯存，常用的cudaMemcpy函數就是通常從CPU拷貝到全局顯存的函數。Global Mem能被所有thread訪問，其在GPU的位置和Cache如下：
- 位置：device memory
- Cache：L1/L2
設備顯存（device memory）
device memory并非位于SM內部，而是由所有SM共享，因此訪問速度較慢，需要Cache緩存加速。除此之外，device memory必須通過32/64/128-byte的【memory transaction】訪問，并且要求這些memory transaction是aligned to their size。
舉例而言，即讀取32-byte的memory transaction時，地址必須是32的倍數；讀取64-byte的mem transaction時，地址必須是64的倍數當一個warp執行指令（load/store）來訪問Global mem時，它會根據【每個thread訪問的word的大小】和【每個thread訪問的地址關系】來把該訪存指令聚合成一個或多個memory transaction。舉例而言，如果每個thread訪問4byte的word，則一個warp（32個thread）就需要訪問32*4=128byte的內存。
- 如果這32個word時連續且對齊的，那么只需要 一個128-byte memory transaction 或 四個32-byte mem transaction即可。
- 如果連續，但起始地址并未對齊128byte，那么需要 兩個128-byte memory transaction 或 五個32-byte mem transaction。
- 如果不連續，那么SM會將能放在一個128-byte mem transaction的thread的訪存操作聚合成一個128-byte mem transaction，因此會產生多個128-byte memory transaction。(32-byte mem trans同理)
針對于每個thread所讀取的word，若word size是1/2/4/8/16 byte，且是Naturally Aligned，則會被編譯成一個memory instruction。（后續同一個warp的memory instruction會進行聚合）如果不滿足size和alignment的條件，那么當前thread的該次mem access就會被編譯為多個mem instruction，因此變慢。
L1/L2 Cache
Global Memory的讀取會被緩存到L2（有時也會緩存到const cache），通過可配置選項可以選擇是否緩存到L1。
- 如果Mem Access同時緩存在L1/L2上，那么是通過128-byte mem transaction來實現的
- 如果Mem Access僅緩存在L2上，那么是通過32-byte mem transaction來實現的
（因此，僅緩存在L2對分散的內存讀取有好處，可以減少over-fetch）
即，L1的Cache Line Size = 128 byte，L2的Cache Line Size = 32 byte。所以當L1/L2共存時，取最大的Cache Line Size。
L2 Cache有如下特點:
- 所有的SM共享一個L2 Cache
- 用來緩存對global/local memory的讀取。
- 有時也會用來處理Register Spilling
（可以通過device property中的l2CacheSize來查看其大小）
局部顯存（Local memory）
每個thread都擁有自己私有的local memory，負責存儲一些局部變量（automatic variable）。
對于局部變量而說，一些小型的局部變量會被放到register里，當register不夠用時，則會被放到Local Mem中。
Local Mem的位置和Cache如下：
- 位置：device memory
- Cache：L1/L2
由于local mem也是放在device memory上，所以其和global mem很像。即access latency和bandwidth都和global mem一樣低，一些內存對齊的約束也得滿足。
有一點local mem獨有的優化是：如果warp內的threads同時訪問相同的local mem里的relative address（e.g. same index in an array variable, same member in a structure variable），memory access are fully coalesced。
共享顯存（Shared Memory）
shared memory位于thread block這一層，即每個block共享一塊shared mem，這塊shared mem對該block內的所有threads可見，且當該block執行結束時，其所占用的shared mem也會被釋放。
shared mem的位置和cache如下：
- 位置：on-chip memory，片上顯存
- cache：無，因為是on-chip的，讀取速度夠快不需要cache
shared mem本身位于芯片上，所以讀取速度很快，可以作為software-managed cache來加速的執行。L1/L2 cache上存儲什么數據無法由程序來直接控制，但我們可以控制shared mem上存儲什么數據。
在硬件上，shared memory 被分成32（對應warp中的thread個數）個相等大小的【bank 內存塊】。
- 每個bank的帶寬是32 bits per clock cycle
- 連續的32-bit words是放在連續的32個banks中
這32個內存塊們可以同時被訪問：
- 若32個thread各自訪問32個bank的word，就只需要一次內存傳輸就行。
- 若不同thread訪問同一個bank的不同32-bit word，就產生了【bank conflict】，access就會被序列化，需要多次內存傳輸。
- 若不同thread訪問同一個bank的相同word，不會產生bank conflict，僅需要一次內存傳輸，此時觸發【broadcast】，word會被廣播給多個thread。
常顯存（Constant Memory）
Constant Memory，顧名思義是用來存儲只讀數據的內存。此處的【只讀】針對的是device code的只讀，我們可以通過Host向Constant Memory寫入數據（通過cudaMemcpyToSymbol()的接口），然后在device code中讀取。常見的Constant Memory大小為64KB，其位置和Cache如下：
- 位置：device memory
- Cache：constant cache（比L1/L2快）
寄存器（Register）
Register位于SM上，每個SM都有固定數目的一組threads。每個thread使用的register越少，就有越多的block/threads可以并發的位于同一SM上，進而提高性能。每個thread使用的register的數目由編譯器【啟發式】的決定。但我們也可以通過Launch Bounds提供一些信息協助編譯器更好的決定。
參考文獻：
CUDA2.2-原理之存儲器訪問 - 仙守 - 博客園
CUDA程序調優指南（一）：GPU硬件 - 知乎
總結
以上是生活随笔為你收集整理的GPU硬件结构和程序具体参数设置的全部內容，希望文章能夠幫你解決所遇到的問題。
- 上一篇： char装cstring_VC中char
- 下一篇： 2018广技师C语言专插本试题,2018