第二章練習題 (Practice - Programming GPUs in CUDA)

答題前先用這張表把每個大節的核心關鍵字與「最常考的陷阱」對上號；看到題目中的關鍵字就能回想對應概念。

用 __global__ 修飾、回傳型別必為 void（結果只能透過指標參數寫入 device 可存取的記憶體）。<<<a, b>>> 第一參數 a 是 grid 維度（block 數），第二參數 b 是 block 維度（每 block thread 數）；1D 可用 integer，多維用 dim3。現行 GPU 每 block 上限 1024 threads，因為一個 block 內所有 thread 都駐留在同一個 SM 並共享其資源。

Unified：cudaMallocManaged 配置 unified memory，單一指標 CPU/GPU 皆可存取、driver 自動搬移。Explicit：cudaMalloc 配 device memory、cudaMemcpy 手動搬移，host pointer 與 device pointer 分離、不可互相 dereference。cudaMemcpy 是同步的（複製完才返回、會阻塞 host thread）。page-locked memory 用 cudaMallocHost 配置（提升複製效能、且是非同步傳輸的必要條件）。兩種 device buffer 都用 cudaFree 釋放，cudaMallocHost 用 cudaFreeHost 釋放。

blocks = ⌈1000/256⌉ = 4（(1000 + 256 - 1) / 256 = 1255 / 256 = 4，或用 cuda::ceil_div(1000, 256)）。4 × 256 = 1024 > 1000，最後 24 個 thread 會越界，故需 bounds check 讓超出範圍的 thread 直接結束、不存取陣列。加上守門後可安全地啟動多於需要的 thread——閒置 thread 開銷小可接受，但應避免啟動整個都閒置的 block。

不能。triple chevron 啟動本身不回傳 cudaError_t；launch 後立即查到 cudaSuccess 只代表 launch 參數/execution configuration 無誤、且當下無殘留錯誤，不代表 kernel 已執行甚至已開始。kernel 與多數 API 為非同步，執行期間的錯誤要等下次檢查狀態才會回報。要抓執行期錯誤須用 cudaDeviceSynchronize() 等待後再檢查（如 CUDA_CHECK(cudaDeviceSynchronize())）。非同步錯誤一旦發生會被之後每個回傳 cudaError_t 的 API 持續回傳，直到 cudaGetLastError 把狀態清除為止。

用 @cuda.jit 裝飾器標記，在第一次 launch 時對當前 GPU 做 JIT 編譯。launch 用方括號 execution configuration：kernel[blocks, threads](args)，對應 C++ 的 <<<blocks, threads>>>（引數順序相同）。cuda.grid(1) 是 cuda.threadIdx.x + cuda.blockIdx.x * cuda.blockDim.x 的 shorthand。本章用 numba.cuda（寫/launch kernel）+ CuPy（GPU 記憶體與 ndarray）；安裝只需最新 NVIDIA Driver，通常不需 CUDA Toolkit。

安全是因為 A/B/C 是 CuPy 建立的 ndarray，array 攜帶各維 extent，會自動做邊界檢查；grid_size 用 ceil 算保證 thread 數 ≥ 元素數，多出的 thread 因自動邊界檢查而安全（對照 C++ 通常需手寫 if (idx < N)）。block_size = 2048 超過任何現行 GPU 的 1024 上限，kernel launch 失敗並拋出 exception（CUDA_ERROR_INVALID_VALUE，參數超出可接受範圍）；Python 用 try/except 捕捉，否則程式異常結束並 dump traceback。

Local memory：scope = thread、lifetime = kernel，但 location 在 device（global）memory space，不在 SM。「local」指的是邏輯 scope（thread-local），不是實體位置；它的延遲/頻寬等同 global memory，不是晶片上快取（是 register spilling 的去處）。實體在 SM 上的：register、shared memory（與 L1 共用實體空間）。實體在 device 上的：global、constant、local、L2。

它是 block 層級 barrier：block 內所有 thread 都到達該呼叫前，沒有任何 thread 能往下執行；並保證 barrier 前對 shared memory 的寫入排序在 barrier 後的讀取之前可見（先寫後讀的可見性）。它只同步單一 thread block 內的 thread，不跨 block（跨 block 需 thread block clusters / cooperative groups / atomics）。若放在只有部分 thread 會進入的分支，未進入的 thread 永遠到不了 barrier，造成 deadlock 或未定義行為。

2D block 中 threadIdx.x 變化最快（出現在 myCol）。讀 a 時 myCol 是 INDX 的第二參數（col），連續 threadIdx.x → 位址 +1（連續）→ 完美合併。寫 c 時 myCol 是 INDX 的第一參數（row），連續 threadIdx.x → 位址 +ld（stride = ld）→ 不合併。當矩陣 > 32（ld > 32）時 stride ≥ 32 floats = 128 bytes ≥ 32 bytes，warp 被迫為每個 thread 各發一個 transaction，利用率僅 12.5%（128 used / 1024 transferred，病態未合併）。解法：用 shared memory 當 user-managed cache 暫存 (stage)，使讀與寫 global memory 皆合併。

會造成 32-way bank conflict。陣列 row-major，線性位址 = row*32 + col，bank = (row*32 + col) % 32 = col % 32。當 threadIdx.x 是第一索引（row）、threadIdx.y 固定時，連續 thread 以 stride 32 元素存取 → 全落在同一個 bank → 序列化。修正：把 column 維度 +1 宣告成 [32][33]，列長變 33（與 32 互質），整欄存取位址 (row*33 + col) % 32 = (row + col) % 32 對不同 row 落在 32 個不同 bank，消除衝突；代價是每列多一個未用的 float。

entry point 用 __tile_global__（類比 __global__），device function 用 __tile__（類比 __device__）；Python 對應 @ct.kernel / @ct.function。triple-chevron 第二參數必須是 1，因為 tile kernel 的 thread 數由 compiler 決定（programmer 視角每個 tile block 由「單一邏輯 thread」執行），寫任何非 1 的值都是錯誤。tile 的 shape/dtype 必須編譯期已知，且每一維必須是 2 的次方。

兩種：(1) tile-space load/store（用 view 物件 + tile-space index，規則、可預測的映射）；(2) gather/scatter（用 index tile / pointer tile，任意非連續位置、逐元素）。tile-space load 可被 compiler 下放到 TMA（Tensor Memory Accelerator），遠快於 per-element gather——能用就別用 gather。C++ 兩步建構：先 ct::tensor_span（raw pointer + ct::extents 賦予多維結構）→ 再 ct::partition_view（切成固定大小 tile 的網格，提供以 tile-space 座標操作的 .load/.store）。

Cross-block contention（各 block 的 partial result 合併進 global memory 同一位置）需 device scope，因結果要對 device 上所有 block 可見（C++ ct::thread_scope_device_t{}）。Intra-block contention（單一 block 內 tile 多元素寫同一位置）只需 block scope（ct::thread_scope_block_t{}）。關鍵差異：C++ thread scope 省略時預設 system-wide，Python 預設 device scope（兩者不同，易混淆）。整個 call 並非 atomic：tile 對每個元素各做一次 atomic update，元素間順序未指定 (unspecified)。要把 tile 加總成 scalar 應改用 tile reduction，而非硬湊 intra-block atomics。

結果形狀為 4x8x2。先 rank promotion：8x2 補 leading singleton → 1x8x2；再 trailing 對齊與 singleton 拉伸（dim0：1→4、dim1：8 vs 1→8、dim2：2=2）→ broadcast shape 4x8x2。GEMM 慣用法：無論輸入精度都以 FP32 累加（accumulator 為 FP32），store 時 cast 成輸出元素型別；K-loop 迭代 ceil(K / tk) 次（= (K + tk - 1) / tk），部分 K-tile 在 load 時 zero-pad，C 側部分 M/N 邊緣 tile 用 store 端 OOB-discard 處理。

stream 是 triple-chevron 的第四個參數：<<<grid, block, shared, stream>>>（不用 shared memory 時第三參數填 0）。非同步傳輸用 cudaMemcpyAsync()（立即返回）；host buffer 必須是 pinned / page-locked（用 cudaMallocHost），否則 cudaMemcpyAsync 會退化成同步、無法重疊。計時用 cudaEventElapsedTime(&ms, start, stop)，單位毫秒 (ms)，取值前須先確保兩個 event 都已觸發（如 cudaStreamSynchronize）；event 用 cudaEventRecord(event, stream) 插入 stream。

cudaStreamCreate 建立的是 blocking stream；legacy default stream（NULL stream / ID 0）也是 blocking 且為所有 host thread 共享，它會與所有 blocking stream 同步——default stream 的 kernel2 會等 stream1 的 kernel1 完成，stream2 的 kernel3 又會等 kernel2 完成，即使三者原本可並行。避免方式：用 cudaStreamCreateWithFlags(..., cudaStreamNonBlocking) 建 non-blocking stream，或啟用 per-thread default stream（--default-stream per-thread 或 CUDA_API_PER_THREAD_DEFAULT_STREAM）；之後不能假設順序，需顯式同步。CUDA Graphs：capture → instantiate（各做一次）→ launch（多次），把重複從 host 發出同一串 API 呼叫的延遲與 CPU 開銷降到最低。

UVA（Unified Virtual Address Space） 讓單一 process 內所有 host memory 與所有 GPU global memory 共用一個虛擬位址空間，CPU 與每個 GPU 各佔獨特範圍；可用 cudaPointerGetAttributes() 由指標判斷其位置，cudaMemcpy 也能用 cudaMemcpyDefault 自動判斷方向。配置 managed memory 三種方式：cudaMallocManaged、cudaMallocFromPoolAsync（managed pool）、帶 __managed__ specifier 的全域變數。在有 HMM 或 ATS 的系統上，所有 system memory 都隱含是 managed memory，無需特別配置。

三個屬性（用 cudaDeviceGetAttribute 查詢）：cudaDevAttrConcurrentManagedAccess（1 = full、0 = limited support）、cudaDevAttrPageableMemoryAccess（1 = 所有 system memory 皆 unified、0 = 只有明確配置的 managed memory）、cudaDevAttrPageableMemoryAccessUsesHostPageTables（1 = hardware coherence / ATS、0 = software coherence / HMM）。cudaHostRegister page-lock 既有記憶體後，不能用 host 指標在 kernel 存取，必須用 cudaHostGetDevicePointer() 取得 device 指標並在 kernel 中使用該指標（host 指標 a 與 device 指標 devA 是兩個不同指標）。

compute_XY 是 virtual architecture，產生可 JIT 的 PTX（向前相容、無 Cubin）；sm_XY 是 real hardware architecture，產生 Cubin（以 SM version 識別），同時也內嵌 PTX 保留向前相容。-arch=native 自動偵測並只為當前 GPU 產生 Cubin，不含 PTX，因而無向前相容。device code 編譯鏈順序：C/C++ device code → PTX → ptxas → Cubin；多個 PTX + Cubin 可內嵌進單一 Fatbin。（nvcc 是 offline 編譯，nvrtc 是 online / JIT。）

需加 -rdc=true（別名 -dc）啟用跨單元 device code linking，稱為 separate compilation（預設是 whole-program compilation，要求所有 GPU code 在同一單元）。它不是預設因為跨檔 device linking 可能影響效能。LTO（Link-Time Optimization） 在 link 時跨各別編譯的檔案做最佳化，挽回大部分 whole-program 的效能、同時保留 separate compilation 的彈性；用 -dlto flag 或 lto_<SM> target 啟用，且 device 編譯（-dc）與最終 link 兩步驟都要帶 LTO 旗標。（陷阱：CUDA 13 起 __global__ 與 __device__/__managed__/__constant__ 變數預設為 internal linkage。）

把全章「一看關鍵字就要反射出的答案」濃縮成下表，臨考前快速掃一遍。