記憶體架構與資料局部性 練習題 (Practice - Memory Access Efficiency and CUDA Memory Types)

ratio = 每從 global memory 存取 1 byte 所執行的 FLOP 數。內層迴圈每次迭代做 1 次 FP multiply + 1 次 FP add = 2 FLOP，並讀 M、N 各 1 個 4-byte 值 = 8 B → ratio = 2/8 = 0.25 OP/B。別名為 arithmetic intensity 或 computational intensity。

吞吐量上限 = bandwidth × intensity = 1555 GB/s × 0.25 FLOP/B = 389 GFLOPS，僅為 FP32 peak 的 2%（對比 tensor core 156,000 GFLOPS 更只有 0.25%）。執行速度被「資料從 memory 送到 core 的速率」卡住，而非運算單元 → 即 memory-bound program。要餵滿 19,500 GFLOPS 需 ratio ≥ 19,500/1555 ≈ 12.5 OP/B。

關鍵陷阱：自動「陣列」放 local memory（實際位於 off-chip global，非 register，除非編譯器發現索引全為常數）。

1. 頻寬/延遲：on-chip register file 聚合頻寬比 global memory 高 ≥2 個數量級；且值存進 register 後該存取不再消耗 global 頻寬，等於提高 compute-to-memory ratio。
2. 指令數：算術指令的 operand 內建為 register（fadd r1, r2, r3）；若 operand 在 global，需先 load 再 fadd，多一條指令多耗 cycle。
3. 能耗：讀 register 的能量比讀 global memory 低 ≥1 個數量級。

(a) 無 tiling：每元素被請求 N 次。(b) T×T tile：協同載入後每元素只被請求 N/T 次 → global memory 流量減量因子 = TILE_WIDTH（不是 T²）。16×16 tile → ratio = 0.25 × 16 = 4 OP/B，A100 吞吐上限提升到 1555 × 4 = 6220 GFLOPS。

• 載入後 (line 21)：防 Read-after-write (RAW) = true dependence（reader 真的需要別 thread 寫入的 tile 資料，非等不可）。
• 用完後 (line 26)：防 Write-after-read (WAR) = false dependence（writer 不需 reader 的資料，僅因重用同一 shared 位置才相依，換位置即可消除）。漏任一個都會造成 race condition。

row-major 線性化下，超出列尾的存取會「靜默」讀到下一列的合法值而汙染內積；超出陣列尾則回傳隨機值或 crash。需 三組獨立檢查：載入 M (Row<Width && ph*TILE_WIDTH+tx<Width)、載入 N (ph*TILE_WIDTH+ty<Width && Col<Width)、寫入 P (Row<Width && Col<Width)，因為三者索引界線不同。越界載入填 0.0f（乘加皆中性，不改變內積）。注意越界發生在所有 phase，不只最後一個。

平均用量 = 32 KB / 256 threads = 128 B/thread，超過 full-occupancy 門檻 164 KB / 2048 = 82 B/thread → 無法 full occupancy。每 SM 可容納 threads ≈ 164 KB / 128 B ≈ 1272–1312（書中算 ≈132 B/thread → 1272），max occupancy ≈ 1272 / 2048 ≈ 62%，限制因子是 shared memory。（對比 tiled matmul 只用 8 B/thread，不受限。）

kernel intensity = 36 FLOP / (7 × 4 B) = 36/28 ≈ 1.29 OP/B。比較機器平衡點 peak FLOPS / peak BW：

• (a) 200/100 = 2.0 OP/B。kernel 1.29 < 2.0 → memory-bound。
• (b) 300/250 = 1.2 OP/B。kernel 1.29 > 1.2 → compute-bound。

同一 kernel 在不同硬體可能落在分界兩側，取決於機器平衡點。

naive：ratio 0.25 → 389 GFLOPS（峰值 2%）。tiled 16×16：協同載入讓每元素只讀 1 次，ratio 升 16 倍至 4 OP/B → 6220 GFLOPS（峰值 32%）。tiling 有效是因為 matmul 具 data reuse / locality，把熱資料放進 on-chip shared memory，移除大量 global 流量。仍非峰值是因為 ratio 4 OP/B 距餵滿所需的 12.5 OP/B 仍有差距（殘留 memory-bound）。進一步改善：thread coarsening、register tiling，或直接用高度優化的 cuBLAS / CUTLASS。

關鍵差異在 scope：register 是 thread 私有，一個 thread 載入的值只有它自己能用；shared memory block 內所有 thread 可見。tiled matmul 的精髓是「協同載入」——一個元素被多個 thread 重複使用（如同 row 的 thread 共用 M、同 col 的 thread 共用 N）。把值放 shared memory，每個 global 元素只需載入一次就能服務多個 thread，達成 TILE_WIDTH 倍的流量減量；若放 register，每個 thread 都得自己重載，無法共享，失去 tiling 的全部好處。

CPU 核心通常一次只跑 1–2 個 thread，重用資料很可能仍留在 cache 中，靠硬體自動管理即可（隱式 blocking）。GPU 的 SM 同時跑大量 thread 以隱藏延遲，這些 thread 會競爭 cache slot，使 cache 留住重用資料變得不可靠。因此 GPU 必須用 shared memory (scratchpad) 顯式、程式員手動管理重要的重用資料，確保它留在 on-chip。兩者目標相同（提升 data locality、把重用資料留在晶片上），但 GPU 因高並行度而需顯式控制。

Strip-mining：把一條長迴圈拆成「外層 phase 迴圈 + 每 phase 執行少數連續迭代的內層迴圈」，並在內層前後加 barrier，強迫同 block 的 thread 每 phase 聚焦於同一段輸入 → 正是 tiling 所需的 phase 機制。phase 數 = Width / TILE_WIDTH（任意寬度時用 ceil(Width/(float)TILE_WIDTH)）。

(a) constant 變數：scope = grid（所有 grid 所有 thread 共用），lifetime = application，容量上限 ≤ 65,536 bytes，device 端 唯讀（存在 global memory 但被 cache，搭配適當存取模式極快）。
(b) 三項：① extern __shared__ 宣告（省略大小、合併成一個一維陣列）；② kernel launch 第三個 <<<grid, block, size>>> 參數傳入位元組數；③ 用 cudaGetDeviceProperties 查 devProp.sharedMemPerBlock 決定 size。需手動切分 Mds/Nds 起點並改用線性化索引。

(a) local 變數 scope = thread → 每個 thread 一份 = 1000 × 512 = 512,000 份。
(b) shared 變數 scope = block → 每個 block 一份 = 1000 份。
重點：版本數由 scope 決定（thread-scope = 總 thread 數；block-scope = block 數；grid-scope = 1 份）。