昨天我們聊的是「什麼叫做優化」。一個 Kernel 跑 100 ms，光看這個數字其實沒有太多意義。你至少得知道原本多快、改了什麼，以及時間究竟花在哪裡。
但要繼續往下談，很快就會碰到一個更基本的問題：GPU 到底在幹嘛？
我第一次接觸 CUDA 的時候，最困擾我的其實不是語法。threadIdx.x、blockIdx.x 看幾次就知道怎麼用了，Warp 是 32 個 Threads 也很好背，Shared Memory 很快、Register 很近，這些句子網路上到處都是。
真正麻煩的是，我不知道它們彼此之間是什麼關係。
Thread 是什麼? 是脆嗎? 怎麼發音?
不是喔 是GPU執行緒 ， 我都念 t h r e a d
Thread 到底在哪裡跑？Block 為什麼存在？Warp 只是把 32 個 Threads 湊在一起嗎？Shared Memory 又為什麼偏偏是 Block 裡的 Threads 可以一起用？
如果這些東西只是一張名詞表，等到真的看到一個 Kernel 很慢的時候，還是不知道該從哪裡看。
所以今天不打算把 GPU Architecture 從頭背一次。我想換一個方式：我們直接走進一顆 GPU 裡面，然後拿最普通的 Matrix Multiplication 當手電筒。
一路看看一個 C[i,j] 是誰算的、它在哪一個 SM 裡跑、旁邊另外 31 個 Threads 又在做什麼，最後再看看為什麼大家寫 GEMM 時，做著做著就會開始使用 Shared Memory。
(請原諒我從這個例子開始 雖然我知道現在大部分不會這樣寫gemm 有tensor ptx等等的 但我覺得matrix multiply是最有幫助我 了解整個GPU系統的一個好的開始 畢竟兩個就很像XD 總之 如果不棄嫌的話 可以隨著我往下看看XD)
如果我們先不負責任的把很多細節拿掉，一顆 GPU 可以先粗略看成這樣：
(真的很不負責任喔 下面的內容也都是 先附上一些我覺得很負責任的來源XD)
一些很可以參考的來源-1
一些很可以參考的來源-2
一些很可以參考的來源-3
GPU
┌────────────────────────────────┐
│ │
│ SM SM SM │
│ SM SM SM ... │
│ │
├────────────────────────────────┤
│ L2 │
├────────────────────────────────┤
│ HBM │
└────────────────────────────────┘
真正執行大量計算的地方，是一個一個的 SM（Streaming Multiprocessor）。
如果把整顆 GPU 想成一間工廠，SM 比較像裡面真正工作的車間。GPU 有很多個 SM，所以同一時間可以攤開很多工作。
一個 SM 裡面，大致會看到這些東西：
┌───────────────────┐
│ SM │
│ │
│ Warp Scheduler │
│ │
│ Execution Units │
│ │
│ Register File │
│ │
│ L1 Cache │
│ │
│ Shared Memory │
└───────────────────┘
Execution Units 是真正做 arithmetic 的地方；Warp Scheduler 負責安排 Warp 的執行；Register File 則提供大量非常靠近 computation 的 thread-local working storage。
Shared Memory 也在 SM 這個局部硬體環境裡，但它的角色跟 Register 不太一樣。Register 通常是 Thread 自己用，Shared Memory 則是一個 Block 裡的 Threads 可以共同使用、而且由 programmer 明確管理的 storage。
更大的 Tensor，例如 A、B、C，則通常還是放在容量大得多的 device memory。GPU 上還有 L1、L2 等 cache 幫忙處理資料存取。
今天不需要先背它們分別有幾 MB、多少 cycle latency。先有一個感覺就夠了：
GPU 有很多地方可以算，也有很多不同距離、不同用途的地方可以放資料。
後面很多 GPU optimization，最後都會回到兩件事情：讓硬體有足夠的工作可以做，以及不要為了同一份資料付出太多不必要的搬運成本。
假設今天要算：
C = A @ B
Matrix Multiplication 裡，每一個輸出元素其實就是一個 dot product：
(對 我也沒有想到我居然需要放這個圖片.. IT邦居然沒有支援latex OMG..)
(我J個甚至是在word上面打完截圖的 所以請原諒我後面有個換行鍵符號 哈哈哈哈)
最直覺的 GPU 寫法，是讓每一個 Thread 負責一個 C[i,j]。
__global__ void matmul(
const float* A,
const float* B,
float* C,
int M,
int N,
int K
) {
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
if (row >= M || col >= N)
return;
float acc = 0.0f;
for (int k = 0; k < K; ++k) {
acc += A[row * K + k] * B[k * N + col];
}
C[row * N + col] = acc;
}
假設某一個 Thread 最後算到：
row = 19
col = 37
那它從頭到尾就負責：
C[19,37]
這就是 Thread 最直接的感覺。
你寫 CUDA Kernel 時，其實是在用「一個 Thread 的視角」寫程式。很多 Threads 都會跑同一份 Kernel code，只是因為自己的 threadIdx、blockIdx 不一樣，所以最後處理的是不同資料。
來 這裡有一個很重要的觀念：Thread 不是一顆小 CPU Core。
我不會考你 但你的直覺可能會唬爛你
你 launch 十萬個 Threads，不代表 GPU 裡真的有十萬顆硬體核心。比較好的理解是：你建立了十萬份工作，接著交給 GPU 去安排它們什麼時候、在哪裡被執行。
CUDA 也不是把所有 Threads 散裝丟進 GPU。Threads 會先被組成 Block，很多 Blocks 再組成一個 Grid。
Grid
│
├── Block 0
│ ├── Thread
│ ├── Thread
│ ├── Thread
│ └── ...
│
├── Block 1
│ └── ...
│
└── Block 2
└── ...
Block 很重要，因為它不只是在幫 Threads 分組。
一個 Block 被安排執行時，會 resident 在某個 SM 上。也因為這群 Threads 在同一個 SM 裡，它們才可以透過 Shared Memory 分享資料，也可以用像 __syncthreads() 這樣的方式彼此同步。
所以我比較喜歡把 Block 理解成：
一群被允許彼此合作的 Threads。
整件事情大概可以想成：
CUDA Program GPU
Grid
│
├── Block 0 ───────────────────────→ SM 2
│
├── Block 1 ───────────────────────→ SM 0
│
├── Block 2 ───────────────────────→ SM 4
│
└── ...
GPU 有很多 SM，所以不同 Blocks 可以被分派到不同 SM 執行。如果 Blocks 比 SM 多很多，也沒有關係，前面的 Block 做完，SM 再繼續接新的工作。
到這裡，我們知道 Grid、Block、Thread 是怎麼把問題切開的。
但進到 SM 裡面之後，還少了一個非常重要的角色：Warp。
假設今天的 Block 是：
blockDim = (32, 8)
也就是一個 Block 有 256 個 Threads。
因為 x dimension 剛好是 32，所以在這個特定例子裡，同一排 Threads 會剛好形成一個 Warp：
Warp 0
Thread (0,0)
Thread (1,0)
Thread (2,0)
...
Thread (31,0)
Warp 1
Thread (0,1)
Thread (1,1)
...
Thread (31,1)
這只是我們刻意挑了一個好理解的配置。Warp 實際上是依 linear thread ID 每 32 個分成一組，不是所有二維 Block 都會「一排剛好一個 Warp」。
Warp 是開始理解 GPU execution 時非常重要的一層。
你寫 Kernel 時，是站在一個 Thread 的視角；但硬體真正執行時，Threads 會以 Warp 為重要的執行群組。Warp Scheduler 會在 SM 裡安排目前可以繼續執行的 Warps。
現在回頭看剛才那個負責 C[19,37] 的 Thread。
它不是自己一個人在工作，你也不是一個人在看這篇文章 學GPU的路上 有妮有窩。
它不是自己一個人在工作，假設它所在的 Warp 正在算：
C[19,32]
C[19,33]
C[19,34]
...
C[19,63]
再看一次 inner loop：
acc += A[row * K + k] * B[k * N + col];
固定 k = 0。
這 32 個 Threads 讀 A 時，大概是：
lane 0 → A[19,0]
lane 1 → A[19,0]
lane 2 → A[19,0]
...
lane 31 → A[19,0]
全部人需要的是同一個 A。
但讀 B 時則是：
lane 0 → B[0,32]
lane 1 → B[0,33]
lane 2 → B[0,34]
...
lane 31 → B[0,63]
是一整排連續資料。
這時候很多原本看起來很抽象的 GPU memory 觀念，就開始有畫面了。
站在一個 Thread 的角度，兩邊都只是「我讀一個 A，再讀一個 B」。但真正到了 Warp 的尺度，Memory System 面對的是這整組 addresses。
所以之後談 GPU memory access 時，不能永遠只問：
這個 Thread 在讀哪裡？
還要開始問：
這一整個 Warp 的 address pattern 長什麼樣子？
這就是 coalescing 之類的概念會出現的原因。
Day 2 先記到這裡就好。到底這些 addresses 最後會變成多少 requests、多少 transactions，留到下一篇再真的拆。(留點扣打給我發 哈哈哈哈 不然我快沒東西講惹)
現在先往另一個更重要的地方看：這些 Threads 其實一直在重複使用同樣的資料。
回到我們負責 C[19,37] 的 Thread。
Kernel 裡有一個：
float acc = 0.0f;
像 acc 這種 Thread 自己反覆使用的小值，在 compiler 與 register pressure 允許的情況下，通常希望保留在 Register 裡。
Register 可以先想成 Thread 執行時很靠近自己的 working storage。
但 A、B 這種大 Matrix 當然不可能全部塞進 Register。它們仍然由更大的 memory system 提供。
問題是，Matrix Multiplication 跑一跑之後，很容易看到一件事情。
同一個 Warp 在算：
C[19,32]
C[19,33]
...
C[19,63]
它們在同一個 k 都會需要：
A[19,k]
而如果把視角拉到整個 Block，不同 output rows 又會重複需要 B 裡同一批資料。
也就是說，GEMM 裡有很多資料其實不是「讀一次、用一次、丟掉」。
它們有很明顯的 reuse。
這時候就會自然想到一件事情：
既然一批資料等等還會被很多 Threads 用到，能不能讓大家合作把它先準備好，然後在靠近 computation 的地方多用幾次？
這就是 Shared Memory 很典型的用途。
A / B
│
cooperative load
│
▼
┌──────────────┐
│Shared Memory │
└──────────────┘
↑ ↑ ↑ ↑
Threads reuse
Shared Memory 可以先想成 Block 的共同工作桌。
一個 Block 裡的 Threads 合作把一小塊 A 和 B 搬進來，等大家都準備好之後，再反覆使用這批資料做 computation。
所以 tiled GEMM 常常會長得像：
// load current tile into shared memory
...
__syncthreads();
// reuse current tile
...
__syncthreads();
// move to the next tile
第一個 barrier 確保資料準備完成後大家才開始用；第二個 barrier 則確保這一批資料不會在還有人使用時，就被下一批資料覆寫。
這時候回頭看 Block，會發現它的意義又更具體了一點。
Block 不只是「很多 Threads 的集合」。
它是一群會在同一個 SM 裡執行、可以共享 Shared Memory、也能彼此同步合作的 Threads。
而這也是為什麼 Shared Memory 的 scope 是 Block。
忠孝東路走九遍(X
到這裡，我們可以重新把 C = A @ B 從外到內跑一次。
我們 launch 一個 Kernel，產生一個 Grid。
Grid 裡有很多 Blocks。
Grid
└── Blocks
Blocks 被陸續安排到 GPU 上不同的 SM。
Block
↓
SM
Block 裡有很多 Threads，而每一個 Thread 可以負責一個 C[i,j]。
Thread
↓
C[i,j]
進入 SM 之後，Threads 會以 32 個一組形成 Warp。
32 Threads
↓
Warp
Warp Scheduler 會安排 SM 裡可以繼續執行的 Warps。
每個 Thread 自己反覆使用的小 working value，例如 acc，通常希望由 compiler 保留在 Register。
Block 裡如果有一批資料想讓很多 Threads 共同使用，可以明確放進 Shared Memory。
而更大的 A、B、C，則由 cache / device-memory system 提供。
當 Warp 執行 memory instruction 時，我們開始關心的不只是某一個 Thread 的 address，而是整個 Warp 的 address pattern。
到了 Matrix Multiplication，又會發現不同 Threads 大量重複使用同一批 A 和 B，所以自然開始考慮 Shared Memory tiling。
把這些角色放在一起，大概是：
GPU
┌───────────────┐
│ SM │
Block ───────────→│ │
│ Warp Scheduler│
│ ↓ │
│ Warps │
│ ↓ │
│ Threads │
│ │
│ Register File │
│ L1 Cache │
│Shared Memory │
└───────────────┘
│
│
L2
│
HBM
而不同資料會依照程式的使用方式，存在 Thread-local 的 Register、Block 共用的 Shared Memory，或由 cache / device-memory system 提供。
到這裡，Thread、Warp、Block、SM、Register、Shared Memory 應該已經不太像六個獨立的名詞了。
它們其實是在描述同一件事情的不同部分：
工作怎麼被切開、被安排到 GPU 上，以及計算需要的資料放在哪裡。
今天故意沒有深入算 transaction，也沒有急著拿一堆 profiler metrics 出來。
我們只看到一件很簡單的事情。
Naïve GEMM 裡，不同 Threads 明明會一直重複使用同一批 A 和 B；Shared Memory tiling 則把這件事情變成一個 programmer 明確安排的合作流程。
直覺上，這應該會比較好。
但到底好多少？原本的 cache 已經幫我們做了多少？一個 Warp 的 32 個 addresses 最後到底會變成多少 memory requests？資料到了 L1、L2、HBM，又分別發生什麼事情？
這些才是下一篇要開始拆的內容。
今天只需要先把 GPU 裡這些角色認清楚。
下一次看到一個 CUDA Kernel，希望腦中至少能跑出這個畫面：
Grid
↓
Block 被安排到 SM
↓
Block 裡的 Threads 被組成 Warps
↓
Warp 在 SM 裡被安排執行
↓
Thread 做自己的 computation
↓
不同資料使用 Register、Shared Memory
或 cache / device-memory system
昨天我們先學會，不要看到 100 ms 就急著談優化。
今天再多學一件事：
Kernel 不是一串 C++ code。
它最後會真的變成一群 Blocks、Warps 和 Threads，在一個一個 SM 裡面跑起來。
當你開始能在腦中看到這個畫面，後面的 GPU performance optimization 才終於有地方可以落地。
如果你想學習更多細節
一些很可以參考的來源:
一些很可以參考的來源-1
一些很可以參考的來源-2
一些很可以參考的來源-3
iThome鐵人賽
看影片追技術