目錄

1、CUDA程序Thread的基本結構

3、單指令多線程架構（SIMT, Single Instruction Multi Thread）

4、硬件多線程（Hardward Multithreading）

5、GPU的顯存結構（Memory Hierarchy）

---

本文主要對GPU的硬件，以及根據硬件定量對參數進行設置，按照先了解硬件在進行參數設置的順序分別進行描述。

1、CUDA程序Thread的基本結構

在邏輯上，threads分為如下三個層次：

thread：每個thread都會運行一次kernel function，threads之間平等無優先級。

block：一組線程，通常放在SM上執行。

grid：一組blocks。通常一次kernel function調用的所有thread都放在一個grid中。

而在硬件上，threads僅有兩個層次：

core：真正執行一個thread的硬件

warp：硬件上并行執行的32個線程，同一warp的thread執行同一條指令

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中的所有threads都是【并發的，concurrently】在一個SM上執行的

多個block的threads也可以并發的在一個SM上執行

如果一個block執行完畢，那么新的block會被放到空出來的SM上執行。

SM的設計目的時并發執行幾百個線程，因此使用了SIMT, Single Instruction Multi Thread?的架構。

在SM內部有兩種級別的并行：

指令級別：單線程（單core）內部的instruction-level parallelism。（芯片內部的指令流水線）。

線程級別：硬件上的多個core的thread-level parallelism（通過硬件上的多線程）

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硬件结构和程序具体参数设置的全部內容，希望文章能夠幫你解決所遇到的問題。

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