CUDA 速查表 (Quick Reference)

一頁式速查：五大主題的關鍵術語、定義、數字與判定條件。每區結尾連回完整概念筆記。

GPU 基礎

術語 / 規則	定義 / 數字
CUDA 推出年份	2006（NVIDIA）；全名 Compute Unified Device Architecture
Pipeline 部分可程式化	2003（僅限繪圖，2006 才能做任意運算且獨立於繪圖 API）
GPU 兩大優勢	更高的 instruction throughput 與 memory bandwidth（前提：相近 price/power envelope）
設計取向	GPU=高吞吐、數千 threads；CPU=低延遲、數十 threads
Transistor 配置	GPU→data processing；CPU→cache + flow control
核心心智模型	GPU 用 throughput 換 latency
上手三途徑	Libraries（cuBLAS/cuFFT/cuDNN/CUTLASS）→ AI frameworks → DSLs（Warp/Triton）
DSL 如何執行	編譯後直接在 CUDA 平台上執行
同樣省電但彈性不如 GPU	FPGA

→ GPU 運算基礎

執行模型與 SIMT

術語 / 規則	定義 / 數字
程式起點	永遠從 CPU (host) 開始
host / device	host=CPU+host memory；device=GPU+device memory
kernel / launch	在 GPU 執行的函式=kernel；launch=平行啟動大量 threads 跑 kernel code
GPU 硬體模型	GPU = 一群 SM，SM 組成 GPC；SM 內含 register file、unified data cache（→shared memory+L1）、functional units
thread 階層	thread → warp → thread block → grid（cluster 為選用層級）
維度	block 與 grid 可 1/2/3 維
同 block 的 threads	在單一 SM 執行，可同步、共用 shared memory
block 間排程	無順序保證，須可任意順序執行，不可有資料相依
cluster 條件	compute capability 9.0+、選用；同一 GPC 同時排程，可用 distributed shared memory（透過 Cooperative Groups）
warp	32 threads 為一組；lane 編號 0–31
SIMT	Single-Instruction Multiple-Threads；每 thread 可走自己控制流、無固定 data-width
warp divergence	同 warp 內走不同分支→未走的 lane 被 mask off；走相同路徑時利用率最高
block thread 數建議	32 的倍數（否則最後一個 warp 有閒置 lane，利用率次佳）

→ 執行模型與 SIMT

Tile 程式設計

術語 / 規則	定義 / 數字
撰寫層級	在整個 thread block 層級寫程式，描述對 tiles 的運算
誰決定 thread 數	compiler（依 tile 運算）；programmer 只給 grid 維度
控制流	block 走單一控制流，支援條件/迴圈，但無 warp divergence 概念
scalar vs tile 運算	scalar 由單一 thread；tile 運算由 block 內所有 thread 平行協作
block vs tile	block=執行單位；tile=資料單位（一個 block 可有多個不同 shape 的 tile）
Array	存於 device memory、mutable、有 shape+dtype、可當 kernel 參數
Tile	block 區域、immutable、各維為 2 的次方且編譯期已知、不一定有記憶體表示、不可當 kernel 參數
Tile space	把 array 切成等大、不重疊 tiles 的索引空間（`⌈M/tm⌉ × ⌈N/tn⌉`）
load vs store 越界	load 依指定處理（如補零）；store 越界靜默丟棄（不對稱）
任意位置存取	gather / scatter
不同 shape 結合	較小者自動 broadcast 擴展
與 SIMT 關係	共存、per-kernel 選擇；同一 tile kernel 可跨架構執行、共用同一 device memory 與硬體

→ Tile 程式設計

記憶體階層

術語 / 規則	定義 / 數字
global memory	device code 視角下 GPU 的 DRAM（GPU 內所有 SM 可存取，非系統各處）
system / host memory	CPU 的 DRAM
Unified virtual memory space	CPU 與所有 GPU 共用單一虛擬位址空間，範圍唯一不重疊→由位址可判斷資料位置
register file	per-SM、存 thread 區域變數、compiler 配置、per-thread
shared memory	per-SM、供 block/cluster 內 threads 交換資料、per-block 配置
register 啟動條件	每 thread register 數 × thread 數 ≤ SM 可用 register；超過則 kernel 無法啟動，須減少 thread 數
L1 cache	per-SM（屬 unified data cache）
L2 cache	較大、全 GPU 所有 SM 共享
constant cache	per-SM、獨立於 L1；快取宣告為 constant 的 global memory，compiler 也可能放入 kernel 參數
explicit 配置	只有對應 device（CPU 或 GPU）能存取，須 cudaMemcpy 明確複製
unified memory	可從 CPU/GPU 存取，runtime/硬體需要時搬移；最佳效能=盡量減少搬移
mapped memory（例外）	GPU 可直接存取的 CPU memory，走 PCIe/NVLINK，延遲高、非高效替代

→ GPU 記憶體階層

CUDA 平台

術語 / 規則	定義 / 數字
Compute Capability	`X.Y`（major.minor）；直接對應 SM 版本，CC 12.0 → `sm_120`
`compute_XY` vs `sm_XY`	`compute_80`=PTX（虛擬 ISA，對應 CC）；`sm_86`=cubin（實體 binary，對應 SM 版本）
NVIDIA Driver	「GPU 的作業系統」，所有 GPU 用途的基礎（含 Vulkan/Direct3D）；版本如 `r580`
CUDA Toolkit	與 driver 分離的獨立產品（libraries/headers/tools）；CUDA runtime 是其中特殊 library（API+語言擴充）
Runtime vs Driver API	runtime API 建構於較低階 driver API 之上；部分功能僅 driver API 提供
PTX	Parallel Thread Execution 虛擬 ISA；可作 IR，版本對應 CC（如 `compute_80`）
cubin	針對特定 SM 版本的實體 binary（如 `sm_120`）
fatbin	容器，可裝多 target cubin + PTX，執行時挑最適 binary
Binary Compatibility	同 major 內 minor `>=` 目標可載入；跨 major 不相容（如 sm_86 可跑 CC 8.6/8.9，不可 8.0/9.0）
PTX Compatibility	PTX 可 JIT 到相同或更高 CC（forward compatibility，不可降階）
JIT	device driver 執行期把 PTX→binary；快取於 compute cache（driver 升級時失效）；缺點=增加載入時間
NVRTC	執行期把 CUDA C++ 編成 PTX（nvcc 的替代）

→ CUDA 平台

第二章：Programming GPUs in CUDA

CUDA C++ 基礎

Kernel 與啟動

項目	速查
Kernel 宣告	`__global__`，回傳型別必為 `void`（結果只能寫指標參數）
啟動方式	triple chevron `<<<grid, block>>>` 或 `cudaLaunchKernelEx`
`<<<a, b>>>`	a = grid 維度（block 數），b = block 維度（thread 數）；4 參數時為 `<<<grid, block, sharedBytes, stream>>>`
每 block thread 上限	1024（同 block 共用同一 SM 資源）
多維	`dim3`，未指定維度預設 1
Index intrinsics	`threadIdx`/`blockIdx`/`blockDim`/`gridDim`，皆 `.x/.y/.z`；Idx zero-indexed
1D global index	`threadIdx.x + blockDim.x * blockIdx.x`
Bounds check	`if(idx < N)`；可啟多餘 threads，避免整 block 閒置
Block 數（ceiling）	`(N + threads - 1) / threads` 或 `cuda::ceil_div(N, threads)`（`<cuda/cmath>`）
Launch 對 host	非同步，立即返回

記憶體管理

配置/釋放	API
Unified（CPU/GPU 單一指標，driver 自動搬移）	`cudaMallocManaged` / `__managed__` → `cudaFree`
Device memory（explicit）	`cudaMalloc` → `cudaFree`
Page-locked host（pinned，async 傳輸必備）	`cudaMallocHost` → `cudaFreeHost`
初始化 device memory	`cudaMemset`

cudaMemcpy(dst, src, bytes, kind)：同步（複製完才返回）。四種 cudaMemcpyKind_t：HostToDevice / DeviceToHost / DeviceToDevice / Default（依指標自動判斷）。
效能準則：避免不必要 host↔device 搬移，資料常駐 GPU、傳輸與運算重疊；unified 調效用 Memory Advise / Prefetch。
例外：ATS/HMM 系統所有 system memory 自動是 unified，無需 cudaMallocManaged。

同步與 Runtime 初始化

cudaDeviceSynchronize()：阻塞 host，等所有先前發出的工作（跨所有 stream）完成；多 stream 改用 cudaStreamSynchronize 或 Events。
完整流程：配置 → H2D copy → launch → sync → D2H copy → 驗證（浮點用 epsilon，禁用 ==）→ 釋放。
Runtime 在第一個需 active context 的 API 時隱式初始化（建 primary context + device code JIT + 載入）。CUDA 12.0+ cudaSetDevice 會初始化 runtime，務必檢查回傳值；cudaDeviceReset 銷毀 primary context；main 之後用 CUDA 介面為 UB。

錯誤檢查與修飾符

API/修飾符	意義
`cudaError_t` / `cudaSuccess`	統一回傳型別；`cudaGetErrorString` 轉文字
`cudaGetLastError` / `cudaPeekAtLastError`	取錯誤狀態並重設 / 取但不重設（per host thread）
`<<<>>>`	不回傳 `cudaError_t`；launch 後立即查狀態
非同步錯誤	須 `cudaDeviceSynchronize()` 才抓執行期錯誤；`cudaGetLastError` 清除狀態
`cudaErrorNotReady`	不算錯誤（來自 stream/event query）
`CUDA_LOG_FILE`	環境變數記錄錯誤（driver r570+）
`__global__`/`__device__`/`__host__`	kernel / GPU 函式 / CPU 函式；`__host__ __device__` 兩端皆可用
變數修飾符	`__device__`→Global、`__constant__`→Constant、`__managed__`→Unified、`__shared__`→Shared
`__CUDA_ARCH__`	偵測 device 編譯路徑分流

Thread Block Cluster：CC 9.0+ 可選階層，cluster 內 block 同排程於同一 GPC，cluster.sync() 硬體同步；portable 上限 8 blocks（小配置會降）；cudaOccupancyMaxPotentialClusterSize 查詢；__cluster_dims__(X,Y,Z)（編譯期）或 cudaLaunchKernelEx（runtime）；用 cluster 時 gridDim 仍以 block 數計。

→ CUDA C++ Kernel 與啟動、CUDA C++ 記憶體管理、同步與完整流程、錯誤檢查與修飾符

CUDA Python

動作	Python（numba.cuda + CuPy）	C++ 對應
標記 kernel	`@cuda.jit`（首次 launch 才 JIT）	`__global__`
launch	`kernel[blocks, threads](args)`	`<<<blocks, threads>>>`
多維	tuple `[(gx,gy),(bx,by)]`；每 block ≤ 1024	`dim3`
index	`cuda.threadIdx/blockDim/blockIdx/gridDim`（`.xyz`）；shorthand `cuda.grid(1)`	`threadIdx`…
配置 device array	`cp.zeros`/`cp.random.`（未指定 dtype 預設 float32*）	`cudaMalloc`
H→D / D→H	`cp.array(host)` / `cp.asnumpy(dev)`（不隱式複製）	`cudaMemcpy` H2D/D2H
device-wide 同步	`cp.synchronize()` / `cuda.synchronize()` / `device.synchronize()`	`cudaDeviceSynchronize()`
錯誤處理	拋 exception，`try/except`；block_size=2048 → `CUDA_ERROR_INVALID_VALUE`	回傳 `cudaError_t`

生態系：控制/跑現成 code 用 cuda.core（≈C++ 的 CUDA Runtime）、cuda.compute、CuPy；寫 kernel 用 cuda.lang(SIMT)、cuda.tile(Tile)。
安裝多在 PyPi（pip），只需最新 NVIDIA Driver，通常不需 CUDA Toolkit。
ndarray 只存 host 或 GPU 其一；傳錯邊 → error；array 攜帶 extent，kernel 內自動邊界檢查（不需手寫 if idx<N）。kernel launch 對 host 非同步。優先用 GPU library，不足才寫 kernel。

→ CUDA Python 入門

SIMT Kernels

Thread 階層與同步

thread → block → grid（皆 1/2/3 維）。gridDim/blockDim 對所有 thread 相同（launch 設定）；blockIdx/threadIdx 各異。
Linearization：x 最快（stride 1）、y stride=blockDim.x、z stride=blockDim.x*blockDim.y；決定 thread→warp 分配（連續 threadIdx.x = 連續 lane）。
__syncthreads()（Python cuda.syncthreads()）：block 層級 barrier，全到齊才放行 + barrier 前寫入有序可見。只同步單一 block；放在 divergent 分支會 deadlock/UB。跨 block 用 clusters / Cooperative Groups / atomics。
warp = 32 threads；warp 內 divergence 傷效能。

裝置記憶體空間（scope / lifetime / location）

記憶體	scope	lifetime	location	重點
Global	Grid	Application	Device	主資料，persistent；kernel 唯一回傳結果途徑
Constant	Grid	Application	Device	read-only，~64KB/device；`cudaMemcpyToSymbol`
Shared	Block	Kernel	SM	user-managed scratchpad，與 L1 共用實體空間
Register	Thread	Kernel	SM	最快，compiler 管理；`-maxrregcount`
Local	Thread	Kernel	Device(global)	register spilling 去處，速度同 global
L2 / L1	—	—	Device / SM	L2 全 SM 共用；L1 與 shared 共用

Shared：static __shared__（編譯期大小）vs dynamic extern __shared__（<<<g,b,bytes>>> 第 3 參數，一個 kernel 僅一個，多個需手動 partition）。
Distributed Shared Memory：CC 9.0+ + cluster + cooperative groups（僅 C++）；總大小 = cluster block 數 × 每 block shared；cluster.map_shared_rank()。
Texture/Surface：現代 GPU 對一般 load 無效能優勢。

記憶體效能

主題	數字/規則
Global memory transaction	32 bytes 為單位
完美合併（連續 thread 取連續 4B word）	128B → 4 個 transaction → 100% 利用率
病態未合併（stride ≥32B）	32×32=1024 bytes 流量、僅用 128B → 12.5% 利用率
合併條件	warp 內存取同批 32B segment（連續或 permuted 皆可），最大化 used/transferred
Shared memory bank	32 banks，連續 32-bit word→連續 bank（word%32），每 bank 32 bits/clock
Bank conflict	同 warp 多 thread 存取同 bank 不同位置 → 序列化；同一位置：讀 broadcast、寫只一個（undefined 哪個）
`smem[32][32]`	整欄存取（stride 32）→ 32-way conflict；整列（stride 1）→ 無
Padding 修正	宣告 `[32][33]`（列長 33 與 32 互質）消除 conflict

Atomics / Cooperative / Occupancy

Atomics：對 global memory 位置 lock→read-modify-write→unlock；cuda::atomic/cuda::atomic_ref（可指定 thread scope，如 thread_scope_device），Python numba.cuda.atomic（add/sub/max/min/compare_and_swap）。節制使用；高效 reduction 用 CCCL（Python cuda.coop）。
Cooperative Groups（C++）：可跨 block / 跨 grid / 跨 GPU 同步，也可定義小於 block 的 group。
Occupancy = active warps ÷ SM 最大 active warps；越高越能隱藏延遲。受限資源：threads/block、shared/block、registers/thread（per-SM 與 per-block 上限）。CC 10.0 範例：<<<512,768>>>→75%、<<<512,32>>>→50%（撞 maxBlocksPerMultiProcessor=32）。查用量 nvcc --resource-usage；限制 register nvcc --maxrregcount（過低會 spill）。

→ SIMT 基礎與 Thread 階層、SIMT 裝置記憶體空間、SIMT 記憶體效能、Atomics/Cooperative/Occupancy

Tile Kernels

結構與啟動

角色	C++	Python（`import cuda.tile as ct`）
Kernel entry	`__tile_global__`	`@ct.kernel`
Device function	`__tile__`	`@ct.function`（可省略）
啟動	`<<<grid, 1>>>`（第二參數必須是 1）/ `cudaLaunchKernel(Ex)`	`ct.launch(stream, grid, kernel, args)`
Block 位置	`ct::bidnum_blocks(`→dim3	`ct.bid(axis)`、`ct.num_blocks(axis)`
編譯期常數	`ct::integral_constant`、`_ic`（`0_ic`）	`ct.Constant[T]`

思維：以**整個 block（tile）**為單位，thread 對應交給 compiler（無 threadIdx、無 warp divergence 概念）。Grid-sizing：(N+tile-1)/tile。
Tile：固定大小多維，shape/dtype 編譯期已知、每維 2 的次方、value semantics（複製成本低，programmer 不配置記憶體）。Factory：zeros/ones/full/iota(C++)/arange(Py)。C++ 型別 ct::tile<T, ct::shape<...>>。共用後端 CUDA Tile IR。

載入/儲存與控制流

方式	定位	邊界
Tile-space load/store	view + tile-space index（規則）→ 可 lower 成 TMA	C++ unmasked partial OOB=UB；`load_masked`(預設填 0)/`store_masked`(靜默丟棄)。Python `padding_mode` ZERO/UNDETERMINED(預設)；store 永遠丟棄 OOB
Gather/Scatter	index/pointer tile（任意）	Python 預設 bounds-safe；C++ 預設不安全，須自建 boolean mask

C++ 建 view 兩步：ct::tensor_spanextents) → ct::partition_view（切 tile，提供 .load/.store）。Python：Array.tiled_view(shape) 或 one-call ct.load(array, index, shape)/ct.store(array, index, tile。
C++ .store(tile, idx) vs Python .store(index, tile)（參數順序相反）。完全在 array 外的 tile = UB（masked 也救不了）。
控制流：每 block 單一控制流路徑，scalar 驅動分支/迴圈；C++ 結構化迴圈 ct::irange+range-for；不允許從迴圈內 return；Python step 須嚴格為正。

運算與基本操作

Broadcasting 遵循 NumPy 語義（scalar 複製、singleton 拉伸、低 rank 對齊 trailing；缺維補 leading singleton）；兩維皆非 singleton 且不等 → ill-formed。型別保留較多資訊者（int+float→float）。scalar 需 narrow：Python promote、C++ ill-formed。
Matmul：matmul=a @ b、mma=a @ b + acc（accumulator 跨 K-tile 累加）。GEMM 慣用 FP32 累加，store 時 cast；K-loop ceil(K/tk) 次；部分 K-tile load 端 zero-pad，部分 M/N store 端 OOB-discard。
Reductions：C++ 永遠保留軸、Python 預設丟掉（keepdims=True 保留）。Scans/cumsum 輸出維度同輸入。Transpose 換前兩軸、permute 任意。Selection：Python ct.where(cond,x,y)、C++ ct::select(cond,lhs,rhs)。Math 函式在 ct namespace（可傳 rounding/subnormal 參數）。compiler 映射 matmul/mma 到 tensor cores。

Atomics 與最佳化

Tile atomic：每元素各一次 atomic、整個 call 非 atomic、元素間順序未指定。Cross-block 用 device scope、intra-block 用 block scope。Python 預設 bounds-check on / ACQ_REL / device scope；C++ scope 省略預設 system-wide（與 Python 不同）。不需舊值用 TiledView.atomic_add（lower 成 PTX atomic reduction）。運算：atomic_and/or/xor/max/min/add/xchng/cas。
Optimization hints：附加 metadata，不改語意、compiler 可忽略、per-construct、可 per-arch。C++ [[cutile::hint(arch, kind=value)]]（arch 用 __CUDA_ARCH__ 慣例 900/1000，0=通用，arch-specific 覆蓋 agnostic）；Python @ct.kernel(...) kwargs / call-site kwargs / ByTarget(...) / replace_hints（autotuning）。Kinds：num_ctas(1,2,4,8,16；sm_80 只 1)、occupancy[1,32]、latency[1,10]、allow_tma（僅 tile-space load/store）。
C++ 效能技巧（TMA 路徑三件套）：__restrict__（保證不重疊，可交錯讀寫）+ ct::assume_aligned(p, 16_ic)（16-byte 對齊，cudaMalloc 保證 ≥16B）+ ct::partition_view（lower 成 TMA）+ ct::irange（pipelining/vectorization）。

→ Tile Kernel 結構、Tile 載入儲存與控制流、Tile 運算與基本操作、Tile Atomics 與最佳化

Asynchronous Execution

Streams 與 Events

API	作用
`cudaStreamCreate` / `cudaStreamDestroy`	建/銷毀 stream（銷毀前先跑完工作）
`<<<grid, block, shared, stream>>>`	在 stream 啟動 kernel（第 4 參數）
`cudaMemcpyAsync(dst,src,size,kind,stream)`	非同步傳輸，立即返回；host buffer 須 pinned 否則退化成同步
`cudaStreamSynchronize` / `cudaStreamQuery`	blocking / non-blocking（空→`cudaSuccess`，未空→`cudaErrorNotReady`）
`cudaEventCreate` / `cudaEventRecord(event, stream)`	建 event / 插入 stream（tracer，追蹤進度、細粒度相依、計時）
`cudaEventSynchronize` / `cudaEventQuery`	blocking / non-blocking
`cudaEventElapsedTime(&ms, start, stop)`	兩 event 夾住量測，單位毫秒 ms

三種同步：blocking（cudaDeviceSynchronize/cudaStreamSynchronize/cudaEventSynchronize）、non-blocking/polling（cudaStreamQuery/cudaEventQuery）、callback。
Stream 內 in-order（FIFO，執行序=enqueue 序）；多 stream 間無順序保證（靠 event/cudaStreamWaitEvent 建跨 stream 依賴）。priority 僅 hint。可重疊：host 運算、device 運算、各方向 memory transfer。

Callbacks / 排序 / Graphs

Host function：cudaLaunchHostFunc(stream, func, data)（cudaStreamAddCallback 將棄用，其 flags 須 0）；callback 內絕不可呼叫任何 CUDA API。
錯誤：常在同步時才浮現；cudaGetLastError(清除)/cudaPeekAtLastError(不清除)/cudaGetErrorName/cudaGetErrorString；CUDA_LAUNCH_BLOCKING=1 定位出錯 kernel。
Default stream：cudaStreamCreate 預設 blocking；legacy default = NULL stream / ID 0，所有 host thread 共享、與 blocking stream 互相同步。non-blocking 用 cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking)。Per-thread default stream（CUDA 7+）：--default-stream per-thread 或 CUDA_API_PER_THREAD_DEFAULT_STREAM。
Priority：cudaStreamCreateWithPriority（數字小=高，預設 0，hint 不搶占）。
CUDA Graphs：capture（cudaStreamBeginCapture/EndCapture）→ instantiate（cudaGraphInstantiate，一次）→ launch（cudaGraphLaunch，多次），降低重複 host-side API 的 CPU 開銷。
提升並行：獨立操作先發、同步盡量延後。

→ 非同步 Streams 與 Events、Callbacks/排序/Graphs

Unified/System Memory

概念	速查
UVA	host + 所有 GPU 共用單一虛擬位址空間；`cudaPointerGetAttributes()` 由指標判位置；`cudaMemcpyDefault` 自動判方向（`cudaDeviceMapHost`/`cudaSetDeviceFlags` 已不需要）
配置 managed	`cudaMallocManaged`、`cudaMallocFromPoolAsync`(managed pool)、`__managed__` 全域變數
判斷模式	`cudaDeviceGetAttribute`：`ConcurrentManagedAccess`(1=full/0=limited)、`PageableMemoryAccess`(1=所有 system memory 皆 unified)、`...UsesHostPageTables`(1=hardware/0=software)
Full（pageable）	所有 system memory（含 `mmap`）皆 unified；first-touch 配置、按需遷移、可 oversubscribe
Hardware coherency（ATS）	Grace Hopper/Blackwell + NVLink C2C；cache line 粒度；native atomics；HMM 自動停用
Software coherency（HMM）	Linux kernel ≥ 6.1.24/6.2.11/6.3；PCIe GPU；page 粒度（`nvidia-smi -q \| grep Addressing`）
Limited（Windows/WSL/Tegra）	只支援明確配置的 managed memory、不可 oversubscribe；GPU 活動時 CPU 不可存取
Advise/Prefetch	`cudaMemAdvise`（放置提示）、`cudaMemPrefetchAsync`（非同步預搬，與 kernel 重疊藏延遲）

Page-locked (pinned) memory：cudaMallocHost（≈cudaHostAlloc 但後者多 flags 如 cudaHostAllocMapped）→ cudaFreeHost；cudaHostRegister 把既有配置 page-lock。async copy 必備、提升 sync copy、可 map 給 GPU。
Mapped memory（zero-copy）：一律 page-locked，資料留 CPU、走 PCIe/NVLink（高延遲低頻寬）。cudaMallocHost/cudaHostAlloc 指標可直接傳 kernel；cudaHostRegister 須 cudaHostGetDevicePointer 取 device 指標。對 mapped memory 的 atomic 對 host/其他 GPU 不保證 atomic。

→ Unified 與 System Memory

NVCC 編譯器

項目	速查
nvcc vs nvrtc	offline 編譯 vs online/JIT runtime compiler
副檔名	`.cu`/`.cuh`=含 device code；`.c/.cpp/.cc/.cxx`=host-only
編譯鏈	device code → PTX → `ptxas` → Cubin；多 PTX+Cubin 內嵌成 Fatbin
`compute_XY` vs `sm_XY`	virtual ISA（PTX，可 JIT 向前相容）vs real hardware ISA（Cubin，以 SM version 識別）
`-arch`	`compute_XY`(PTX only)、`sm_XY`(PTX+Cubin)、`native`(僅當前 GPU、無 PTX/無向前相容)、`all`/`all-major`；`-gencode` 多架構
預設架構	支援的最低架構（最大相容）
host 編譯	`-ccbin`/`NVCC_CCBIN` 指定 host compiler、`-Xcompiler` 透傳；預設靜態 `libcudart_static`，`--cudart=shared` 改動態
Separate compilation	`-rdc=true`/`-dc` 啟用跨單元 device linking；const `__device__` 須 `extern`；CUDA 13 起 `__global__`/`__device__`/`__managed__`/`__constant__` 預設 internal linkage
LTO	`-dlto` 或 `lto_<SM>` target（compile 與 link 都要帶），挽回 separate compilation 效能
Debug	`-g`(host)、`-G`(device，定義 `__CUDACC_DEBUG__`、抑制最佳化)、`-lineinfo`(行號，不影響效能)；GPU code 預設 -O3
Optimization 輸出	`-Xptxas=-maxrregcount=N`、`-res-usage`、`-Xptxas=-warn-spills`/`-warn-lmem-usage`
Profiling	Nsight Compute/Systems；`-lineinfo`、`-src-in-ptx`(需 `-lineinfo`)
Fatbin 壓縮	預設壓縮（`--compress-mode` 預設 `speed`；另 size/balance/none；`-no-compress` 全關）
編譯加速	`-t N`(多架構平行)、`-split-compile N`(最佳化階段平行)、`-time`(CSV 耗時)
其他	`-v`(顯示流程)、`-keep`/`--keep-dir`(保留中間檔)；`-std=c++03..c++23`、`-extended-lambda`、`-expt-relaxed-constexpr`

→ NVCC 編譯器

第三章：Advanced CUDA

進階 API、進階 Kernel 程式設計、Driver API、Multi-GPU 與功能導覽的關鍵術語、API、數字與判定條件。每區結尾連回完整概念筆記。

進階 API（Launch / Clusters / Dependent Launch / Batched Transfer / Env Vars）

進階啟動與 Clusters

項目	速查
triple chevron 四參數	block 維度 / grid 維度 / dynamic shared memory（預設 0）/ stream（預設 default stream）
`cudaLaunchKernelEx`	以 `cudaLaunchConfig_t`（`gridDim`/`blockDim`/`dynamicSmemBytes`/`stream` + `attrs`/`numAttrs`）設定，附加零或多個 `cudaLaunchAttribute`；不改 kernel 原始碼即可加屬性/提示
常用 attribute	`cudaLaunchAttributePreferredSharedMemoryCarveout`（L1/shared 平衡）、`cudaLaunchAttributeClusterDimension`（cluster 大小）、`cudaLaunchAttributeProgrammaticStreamSerialization`（PDL）
Thread block clusters	CC 9.0+ 選用層級；保證同 cluster 的 blocks 同時在單一 GPC 執行，可跨 block 交換資料/同步；portable 上限 8 blocks（小配置會降，`cudaOccupancyMaxPotentialClusterSize` 查詢）
`cudaLaunchAttributeClusterDimension`	CC 9.0+，指定必需（required） cluster 維度（`clusterDim` 3 維）；grid 各維須可被 cluster 維度整除；可 runtime 逐次改變（等效編譯期 `__cluster_dims__`）
`cudaLaunchAttributePreferredClusterDimension`	CC 10.0+，指定偏好（preferred） cluster 維度；須為最小維度整數倍；不保證所有 cluster 採用 → kernel 須在最小/偏好維度皆正確
`__cluster_dims__((x,y,z))`	編譯期固定 cluster 維度；`<<<>>>` 第一參數仍是 block 數，cluster 數由 grid/cluster 隱式算出
`__block_size__((block),(cluster))`	編譯期同時定 block+cluster；啟用後 `<<<>>>` 第一參數變 cluster 數（Blocks as Clusters）；第二 tuple 未給預設 `(1,1,1)`
Blocks as Clusters 陷阱	`__block_size__` 第二 tuple 與 `__cluster_dims__` 不可同時、不可配空 `__cluster_dims__`；帶 smem/stream 時 `<<<>>>` 第二參數須佔位符 1，否則 UB

→ 進階啟動與 Clusters

進階 Streams、同步與 PDL

項目	速查
stream 序列化	同一 stream 內預設序列化；唯一例外 = PDL。跨 stream 並行需：無 event 相依、無 implicit sync、資源足夠
NULL stream 阻斷	中間對 NULL stream 下任何指令會阻斷跨 stream 並行；用 `cudaStreamCreateWithFlags(..., cudaStreamNonBlocking)` 避免（建議一律 non-blocking）
最小同步原則	選剛好夠用、最不一般化者：`cudaStreamSynchronize` 勝 `cudaDeviceSynchronize`；非阻塞用 `cudaStreamQuery`/`cudaEventQuery` polling
四種顯式同步	`cudaDeviceSynchronize`（全 device）/ `cudaStreamSynchronize`（單 stream）/ `cudaStreamWaitEvent`（跨 stream 相依）/ `cudaStreamQuery`（非阻塞查詢）
跨 stream 相依	`cudaStreamWaitEvent()` + non-timing event；只表達相依的 event 用 `cudaEventCreateWithFlags(..., cudaEventDisableTiming)` 提速
未 record 的 event	查詢/等待永遠回傳 success（須自行確保已 `cudaEventRecord`，否則隱性 bug）
Implicit sync 觸發	pinned host alloc / device alloc / memset / 同 device 兩位址 memcpy / NULL stream 任何指令 / L1↔shared 配置切換
提升並行兩準則	獨立操作先發、同步盡量延後
Stream priority	`cudaStreamCreateWithPriority`；範圍 `cudaDeviceGetStreamPriorityRange` 回 [greatest, least]；僅 hint、不搶佔執行中工作、不重評估 work queue
PDL 三要件（缺一不可）	primary `cudaTriggerProgrammaticLaunchCompletion()` + secondary `cudaGridDependencySynchronize()` + secondary launch attr `programmaticStreamSerializationAllowed = 1`
PDL 啟動方式	primary 用一般 `<<<>>>`；secondary 用 `cudaLaunchKernelEx` + config；可覆蓋 secondary 的 launch overhead 與部分執行；重疊是「機會」非保證

→ 進階 Streams 與相依啟動

批次傳輸與環境變數

項目	速查
Batching 目的	把多個（通常較小）task 合成單一較大 operation，攤提個別 dispatch 的 CPU/driver overhead（與 CUDA Graphs、PDL 同理）
批次傳輸 API	`cudaMemcpyBatchAsync` / 3D 變體 `cudaMemcpyBatch3DAsync`；傳 src/dst/size 陣列 + `cudaMemcpyAttributes` + `attrsIdxs`
`attrsIdxs[i]`	第 i 個 attribute 套用的「第一個 transfer index」；`failIdx = nullptr` 安全（只是不記錄失敗 index）
`srcAccessOrder`	`...Stream`（pinned/managed，block 到先前 kernel 完成）/ `...DuringApiCall`（ephemeral stack 指標）/ `...Any`（heap 非 ephemeral 且無 hw managed/coherent access → 立即 stage）
Location hints	`srcLocHint`/`dstLocHint` 用 `cudaMemLocation`（type + id）；與 `cudaMemPrefetchAsync` 共用同一結構
SM vs CE flag	`cudaMemcpyFlagPreferOverlapWithCompute` 偏好 CE 並與 compute overlap；非 Tegra 平台忽略。SM=快但佔算力；CE=慢但釋放 SM、整體 app 可能更快（僅 Tegra 可選）
`CUDA_DEVICE_MAX_CONNECTIONS`	增大可減少跨 stream 獨立工作因 false dependency 被序列化；MPS 下預設較低
`CUDA_MODULE_LOADING`	預設 lazy；設 `EAGER` 把 module 載入移到初始化階段（latency-sensitive 較佳）；lazy 下可加 "warm-up" 呼叫近似 eager
設環境變數時機	啟動 application 前設定；app 內設定可能完全無效

→ 批次傳輸與環境變數

進階 Kernel 程式設計（PTX / 硬體模型 / ITS / Thread Scopes / Async Barriers / Pipelines / Async Copies / L1-Shared）

使用 PTX 與硬體模型 / SIMT / Independent Thread Scheduling

項目	速查
用 PTX 兩種方式	`cuda::ptx` namespace（libcu++）/ inline PTX（`asm volatile`）；最後手段，只用於極度效能敏感處
inline PTX	`asm volatile("add.s32 %0,%1,%2;" : "=r"(r) : "r"(a),"r"(b))`；`volatile` 防編譯器最佳化掉/重排
SM 執行特性	SIMT、in-order issue、不做 branch prediction / speculative execution；little-endian；ILP + 硬體多執行緒 TLP
warp	32 threads；half-warp=16；quarter-warp=8；一次執行一條共同指令，32 路同徑時全效率
branch divergence	只發生在 warp 內，各被取路徑序列化執行、停用不在該路徑 threads；跨 warp 互不影響
SIMD vs SIMT	SIMD 把 width 暴露給軟體；SIMT 指令描述單一 thread 的執行與 branch 行為
Independent Thread Scheduling	CC 7.0+（Volta）；每 thread 獨立 program counter + call stack、可 sub-warp 粒度 diverge/reconverge；不再保證 warp lockstep
修正 warp-synchronous code	顯式 `__syncwarp()`，不可依賴隱式 lockstep（舊免同步 intra-warp reduction 須重檢）
同址寫入	非 atomic：序列化次數依 CC、最終寫入者 undefined；atomic RMW：全部序列化但順序 undefined
Hardware multithreading	warp context 常駐 on-chip → 切換 zero-overhead；scheduler 每 cycle 挑 ready warp 發射；`warps_per_block = ceil(T/32)`
資源不足	連一個 block 都放不下 → kernel launch 失敗
裝置端非同步引入	CC 8.0（Ampere） async copy + async barriers；CC 9.0（Hopper） TMA + async transaction barriers + async MMA
async operation	由 CUDA thread 發起、由 async thread 執行；用 synchronization object（barrier 或 pipeline）signal 完成
Proxy	normal load/store → generic proxy；`LDGSTS`/`STAS`/`REDAS` → async thread on generic proxy；TMA/`tcgen05.`/`wgmma.mma_async.` → async thread on async proxy；跨 proxy 需 proxy fence

→ 使用 PTX 與硬體模型

Thread Scopes 與 Scoped Atomics

項目	速查
thread scope 定義	哪些 thread 能觀察某 thread 的 load/store + 哪些能用 atomics/barriers 互相同步；每 scope 對應一個 point of coherency
五級 scope（窄→寬）	`thread`（–, 無）→ `block`（`.cta`, L1）→ `cluster`（`.cluster`, L2）→ `device`（`.gpu`, L2）→ `system`（`.sys`, L2 + connected caches）
暴露介面	CUDA PTX（`.cta`/`.cluster`/`.gpu`/`.sys`）+ libcu++（`cuda::thread_scope_*`）
scoped atomics 兩組成	Thread Scope（誰可觀察效果）+ Memory Ordering（相對其他記憶體操作的順序約束）
兩種 API	libcu++ `cuda::atomic`/`cuda::atomic_ref`；built-in `__nv_atomic_`（`__NV_ATOMIC_` + `__NV_THREAD_SCOPE_*`）
memory order	`relaxed`（只保原子性）/ `acquire`（load）/ `release`（store）/ `acq_rel`（RMW: fetch_add/CAS）/ `seq_cst`（最強，全域總順序）
單純計數	`memory_order_relaxed`（只需原子性，順序不影響正確性）
producer-consumer	producer `release` store + consumer `acquire` load 形成 happens-before；兩端 relaxed 會資料競爭
效能三原則	scope 最窄、ordering 最弱、位置最近（shared > global atomics）；但 scope 必須涵蓋所有需互相同步的 thread，不可為快選錯

→ Thread Scopes 與 Scoped Atomics

非同步 Barriers 與 Pipelines

項目	速查
async barrier	把 arrive 與 wait 分離（vs 單階段 `__syncthreads()`）；arrive 後可做獨立工作重疊等待時間
可用性	CC 7.0+；shared-memory 硬體加速需 8.0+，可同步 block 內任意 subset（先前只 whole-warp `__syncwarp()` / whole-block `__syncthreads()`）
API	`cuda::barrier`（libcu++ ISO `std::barrier` + CUDA scope 擴充）、低階 `cuda::ptx`、shared-memory primitives；`cuda::device::barrier_native_handle()` 取 native handle
使用流程	`init(&bar, count)` 設 expected arrival count（通常 `block.size()`）→ `token = bar.arrive()`（不阻擋，含 seq_cst/block fence）→ 獨立工作 → `bar.waitmove(token)`（阻擋）
arrival_token	標記目前 barrier phase；phase 完成自動進下一 phase 並重設計數，可迴圈重用
scope 限制	cluster remote shared memory barrier 只允許 arrive、不允許 wait；device/system scope 無硬體加速且只 `cuda::barrier` 支援
低階 wait	`cuda::ptx::mbarrier_try_wait` / `__mbarrier_try_wait`（第三參數逾時 ns）迴圈自旋
pipeline	多階段 FIFO deque；producer 進 head、consumer 出 tail；double/multi-buffering，可預取多個 buffer 重疊 copy 與運算
`cuda::pipeline` API	`producer_acquire` → `producer_commit` / `consumer_wait`（等最舊 stage）→ `consumer_release`（釋放回 pipeline 供再 acquire）
primitives API	`__pipeline_memcpy_async`（global→shared）/ `__pipeline_commit` / `__pipeline_wait_prior(N)`（等除最後 N 次 commit 外全部）；等價 `thread_scope_thread`，有 size/alignment 限制

→ 非同步 Barriers 與 Pipelines

非同步資料複製與 L1/Shared 配置

項目	速查
async copy 定位	單一 kernel 內部 global↔on-SM memory 搬移，解耦「發起」與「完成」；非前章 CPU↔GPU/`cudaMemcpyAsync`
同步 copy 代價	`shared[i]=global[j]` 編譯為 global→register→shared（經 register file 中轉），須 copy 後 + compute 後各一次 `block.sync()`
`memcpy_async`	`cooperative_groups::memcpy_async(block,dst,src,n)` + `wait(block)`；彷彿由另一 thread 執行；wait 前讀寫 = data race
API 層級	Cooperative Groups / libcu++ `cuda::memcpy_async` / `cuda::ptx` / primitives
硬體機制	LDGSTS（8.0+, global→`shared::cta`, 小規模）/ TMA（9.0+, 大型多維 bulk, global↔shared cta/cluster）/ STAS（9.0+, registers→`shared::cluster`）
Unified data cache	L1 與 shared 共用同一物理資源，可在 per-kernel 基礎調切分比例
設 carveout	`cudaFuncSetAttribute(kernel, cudaFuncAttributePreferredSharedMemoryCarveout, v)`；整數百分比或 enum `Default`/`MaxL1`/`MaxShared`；函式須 `__global__`
carveout 是 hint	driver 可改；百分比無法對齊支援值時取下一更大（cc 12.0：100KB 最大、50% → 64KB 而非 50KB）
vs `cudaFuncSetCacheConfig`	後者為硬性需求，交錯不同 shared 設定會序列化 launch（thrashing）；偏好 `cudaFuncSetAttribute`
>48KB shared	須 dynamic shared（`extern __shared__`）+ `cudaFuncAttributeMaxDynamicSharedMemorySize` opt-in，並在 `<<<g,b,bytes>>>` 第三參數傳大小；架構特定

→ 非同步資料複製與 L1/Shared 配置

CUDA Driver API

項目	速查
Driver API 定位	runtime API 建構於其上；入口點前綴 `cu`（runtime=`cuda`）；handle-based、imperative；部分介面（VMM）只在 driver API 暴露
初始化	任何 driver API 呼叫前必先 `cuInit()`，再建立並 current 一個 context
Handle 型別	`CUdevice`/`CUcontext`/`CUmodule`/`CUfunction`/`CUdeviceptr`/`CUarray`/`CUstream`/`CUevent` 等不透明 handle
典型流程	`cuInit`→`cuDeviceGet`→`cuCtxCreate`→`cuModuleLoad`→`cuMemAlloc`/`cuMemcpyHtoD`→`cuModuleGetFunction`→`cuLaunchKernel`
Context	類比 CPU process；各有獨立 address space（不同 context 的 `CUdeviceptr` 指不同記憶體）；host thread 同時 1 個 current，維護 context stack
context stack	`cuCtxCreate`/`cuCtxPushCurrent` push、`cuCtxPopCurrent` detach（還原前一個 current）；無 current 呼叫 → `CUDA_ERROR_INVALID_CONTEXT`
usage count	`cuCtxCreate`=1、`cuCtxAttach`++、`cuCtxDetach`--；歸 0 或 `cuCtxDestroy` 即銷毀
primary context	runtime 隱式建立、per-device 共享；driver 端 `cuDevicePrimaryCtxRetain()` 取得（≠ `cuCtxCreate` 自建 context）
Module	類比 DLL，符號在 module scope；`cuModuleLoad`（檔案）/ `cuModuleLoadData(Ex)`（記憶體影像）載 cubin/PTX/fatbin；`cuModuleGetFunction`/`cuModuleGetGlobal` 取符號
PTX vs binary	跑未來架構須載 PTX（載入時 driver JIT）；binary 架構特定不相容未來；JIT 錯誤用 `CU_JIT_ERROR_LOG_BUFFER`
多 PTX 連結	`cuLinkCreate`→`cuLinkAddData`（`CU_JIT_INPUT_PTX`）→`cuLinkComplete`→`cuModuleLoadData`→`cuLinkDestroy`
Kernel launch	`cuLaunchKernel`（grid/block 各 3 維 + sharedMem + stream + args + extra）；傳參用 pointer 陣列或 `CU_LAUNCH_PARAM_BUFFER_POINTER`（`CU_LAUNCH_PARAM_END` 結尾）
對齊規則	`float4`=16、`float2`=8、`CUdeviceptr`=`__alignof(void*)`；device 端 `double`/`long long` 永遠 two-word；struct 對齊 = 各 field 最大值
Runtime/Driver 互通	driver 建 context 後 runtime 沿用、不另建；`cuCtxGetCurrent` 取 runtime 的 context；`CUdeviceptr` ⇄ `void`/`float` 可 cast；driver 程式可呼叫 cuFFT/cuBLAS

→ CUDA Driver API

Multi-GPU 程式設計

項目	速查
多 GPU 支柱	host context 管理 / UVA / P2P bulk 傳輸 / fine-grained P2P load-store / 上層抽象（IPC、NCCL、NVSHMEM、GPUDirect RDMA）
Device enumeration	`cudaGetDeviceCount()`；`cudaGetDeviceProperties()` 填 `cudaDeviceProp`（`major`/`minor` = compute capability）
current device	`cudaSetDevice()`（任意時刻可切）；alloc/launch 落在 current device；首次呼叫前預設 device 0
stream/event 綁定	建立時綁 current device；kernel launch 到非 current device 的 stream → 失敗；memory copy → 成功
跨裝置 event	`cudaEventRecord`/`cudaEventElapsedTime` 跨裝置 → 失敗；`cudaEventSynchronize`/`Query`、`cudaStreamWaitEvent` 跨裝置 → 成功（後者可做跨裝置同步）
default stream	各裝置各有 default stream；跨裝置 default stream 間無順序保證（可亂序/並行）
P2P 傳輸	`cudaMemcpy`（`DeviceToDevice`/`Default`）或 `cudaMemcpyPeer`/`Async`/`cudaMemcpy3DPeer`（指定 src/dst 裝置）；啟用 P2P 後不經 host 中轉、走 copy engine + NVLink；NULL stream 跨裝置 copy 有同步語意
P2P access	`cudaDeviceCanAccessPeer()` 查詢、`cudaDeviceEnablePeerAccess()` 啟用；UVA 同一指標定址兩裝置，kernel 可直接 deref 對方記憶體
P2P 上限/成本	非 NVSwitch 系統每裝置 peer 連線上限 8；`EnablePeerAccess` 對 peer 所有分配全域生效、隨 peer 數呈乘性開銷 → 改用 VMM API 按需標 peer-accessible
P2P 一致性	跨裝置同步屬 `thread_scope_system`；peer device memory atomic RMW 僅限單一 GPU 存取該物件時保證
managed memory	多 GPU 需系統具 P2P 支援
IOMMU / PCI ACS	Linux bare-metal 須關閉 IOMMU（否則 silent memory corruption）；VM pass-through 開 IOMMU + VFIO；Windows 無此限；PCI ACS 把流量繞經 CPU root complex → 降速
collective	CUDA 不提供 multi-GPU collective；由 NCCL / NVSHMEM 提供

→ 多 GPU 程式設計

CUDA 功能導覽

類別	功能 / 速查
3.5.1 Improving Kernel Performance	async barriers（§4.9）/ async data copies + TMA（§4.11）/ pipelines（§4.10）/ work stealing + cluster launch control（§4.12）；專注 kernel 內部效能
Work stealing	靠 cluster launch control，CC 10.0（Blackwell）；block 可 cancel 尚未啟動的 block/cluster → 奪取其 index → 立即執行，達細粒度 load balancing
3.5.2 Improving Latencies	green contexts（§4.6）/ stream-ordered alloc（§4.3）/ CUDA graphs（§4.2）/ PDL（§4.5）/ lazy loading（§4.7）；不含 kernel 內 memory access latency
Green contexts	= execution context；限 kernel 只用部分 SM，保留的 SM 其他 context（含 primary）不佔用；runtime 自 CUDA 13.1 起支援
stream-ordered alloc	`cudaMallocAsync`/`cudaFreeAsync` 排入 stream、依序生效（vs `cudaMalloc`/`cudaFree` 立即執行）
CUDA graphs	由 stream capture 或 graphs API 建；兩大效益 = 降 CPU launch 成本 + 「整 workload 已知」的最佳化；適合重複 workload（可橫跨多類）
3.5.3 Functionality	EGM（§4.17，需 NVLink-C2C，存取系統內所有記憶體）/ dynamic parallelism（§4.18，從 GPU kernel 發起新 kernel）
3.5.4 Interoperability	與 Direct3D / Vulkan 共享 GPU buffer（§4.19）/ CUDA IPC（§4.15，跨 host process 共享 GPU buffer）；同一機制亦用於多節點 GPU-to-GPU
3.5.5 Fine-Grained Control	VMM（§4.16，經 driver API 控 UVA 佈局）/ driver entry point access（§4.20，自 CUDA 11.3 取 Driver/Runtime 函式指標）/ error log mgmt（§4.8，環境變數 `CUDA_LOG_FILE` + error callback）

→ CUDA 功能導覽

第四章：CUDA Features

Unified Memory（統一記憶體）

完整支援與一致性

項目	速查
兩類系統	hardware-coherent（Grace Hopper，CPU/GPU 共用合併 page table，cache-line 粒度）vs software-coherent（HMM 等，各自 page table，靠 page fault + migration，page 粒度）
HMM 需求	Linux kernel 6.1.24+/6.2.11+/6.3+、CC 7.5+、driver 535+、Open Kernel Modules
完整支援能力	device 可存取 host 任意記憶體（`malloc`/stack/static/global/`extern`/file-backed `mmap`）；global 變數須以指標傳入 kernel（預設 `__host__`-only）
關鍵 device 屬性	`cudaDevAttrDirectManagedMemAccessFromHost=1`（host 無 fault 直讀 GPU 記憶體）、`cudaDevAttrHostNativeAtomicSupported=1`（CPU 記憶體硬體 atomic）、`cudaDevAttrConcurrentManagedAccess`、`pageableMemoryAccess`
managed 直存提示	`cudaMallocManaged` 要 host direct access 須 `cudaMemAdviseSetAccessedBy` + `cudaMemLocationTypeHost`（system-allocated 不需）
Atomic 陷阱	無 `hostNativeAtomic`（含 HMM）→ file-backed atomic 不支援；software-coherent 對 file-backed 做 device atomic = UB（僅 hardware-coherent 合法）
Page size	GPU 偏好 2MiB；GPU TLB miss 比 CPU 貴；大頁→碎片多、TLB miss 少、migration 貴
IPC / 其他	CUDA IPC 不支援 managed memory，但 system-allocated 具 IPC；access counter migration 僅 hardware-coherent（CUDA 12.4+ 支援 system-allocated）

平台限制與效能提示

項目	速查
only-managed 裝置	CC 6.x+ 無 pageable access：managed 完整 coherent，但 GPU 不能存取 system-allocated memory
Windows/WSL/Tegra	CC < 6.0 或 `concurrentManagedAccess=0`：無 on-demand 細粒度遷移、不可 oversubscribe、CPU/GPU 不可並行存取（kernel 跑時 CPU 碰 managed → segfault）；缺 GPU page faulting
Multi-GPU managed	home = active device，他者經 PCIe 降頻；Linux 用無 P2P GPU → 全遷回 system memory；`CUDA_VISIBLE_DEVICES`、`CUDA_MANAGED_FORCE_DEVICE_ALLOC`
stream 綁定	`cudaStreamAttachMemAsync`（獨佔縮為 per-stream）、`cudaMemAttachHost`（初始對 device 不可見）
效能提示（不影響正確性）	`cudaMemPrefetchAsync`（async stream-ordered）、`cudaMemAdvise`（`SetReadMostly`/`SetPreferredLocation`/`SetAccessedBy`）、`cudaMemDiscardBatchAsync`、查詢 `cudaMemRangeGetAttribute(s)`
ReadMostly + 多 GPU prefetch	= read duplication（複本，非遷移）；oversubscription 需 GPU page faulting

→ Unified Memory：完整支援、Unified Memory：平台與效能提示

CUDA Graphs

結構與擷取

項目	速查
三階段	definition → instantiation（`cudaGraphInstantiate`→`cudaGraphExec_t`）→ execution（`cudaGraphLaunch`）；instantiate 一次、launch 多次
節點類型（12 種）	kernel / CPU host func / memcpy / memset / empty / event wait / event record / external semaphore signal / external semaphore wait / conditional / memory / child graph
兩種建法	顯式 `cudaGraphCreate`+`cudaGraphAddNode`；或 stream capture `cudaStreamBeginCapture`/`cudaStreamEndCapture`（不可用 `cudaStreamLegacy`/NULL；可用 `cudaStreamPerThread`；`cudaStreamBeginCaptureToGraph`）
Edge data（12.3）	outgoing/incoming port + type；零初始化=完整依賴；唯一非預設用途 = PDL（`cudaGraphDependencyTypeProgrammatic`）；僅 kernel node 有額外 outgoing port
執行緒安全	`cudaGraph_t` 非執行緒安全；`cudaGraphExec_t` 不能與自身並行

更新與條件節點

項目	速查
何時可 update	拓撲/節點類型不變、只改參數；更新在下次 launch 生效
更新 API	whole：`cudaGraphExecUpdate`（拓撲須相同）；individual：`cudaGraphExec*NodeSetParams`（跳過拓撲檢查）；`cudaGraphNodeSetEnabled`（僅 kernel/memset/memcpy，停用=空節點）
限制	memset/memcpy 只有 1D 可改；kernel 不能改 owning context、不能改是否用 dynamic parallelism
Conditional	IF（size1；size2 含 else，cond==0 執行）/ WHILE（每次 body 後重評）/ SWITCH（執行第 n 個 body，越界不執行）；condition 在 device 評估
Conditional handle	`cudaGraphConditionalHandle`（先於節點建立、關聯單一節點、無法 destroy）；device 端 `cudaGraphSetConditional`；`cudaGraphCondAssignDefault`（否則初值 undefined）

記憶體節點與裝置端啟動

項目	速查
Memory node	`cudaGraphNodeTypeMemAlloc`/`cudaGraphNodeTypeMemFree`；VA 在整個 graph 生命固定（底層 physical 變動不需 update）；GPU ordered lifetime
重用/釋放	圖內共用 VA（指標可能不唯一），圖間 virtual aliasing；graph 銷毀不自動釋放活配置；`cudaGraphInstantiateFlagAutoFreeOnLaunch`
Footprint	`cudaDeviceGraphMemTrim`；`cudaGraphMemAttrReservedMemCurrent`/`UsedMemCurrent`（`cudaDeviceGetGraphMemAttribute`）；`cudaGraphUpload` 把映射成本與 launch 分離
Device graph launch	`cudaGraphInstantiateFlagDeviceLaunch`（僅 host 端 instantiate/update）；3 模式 = fire-and-forget（上限 120）/ tail launch（pending 上限 255）/ sibling；`cudaGetCurrentGraphExec`
同時 launch	裝置端重複 launch 同 graph → `cudaErrorInvalidValue`；host+device 同時 → UB

→ CUDA Graphs：結構與擷取、CUDA Graphs：更新與條件節點、CUDA Graphs：記憶體節點與裝置端啟動

記憶體配置與群組（Stream-Ordered Allocator / Cooperative Groups / PDL）

Stream-Ordered Memory Allocator

項目	速查
動機	`cudaMalloc`/`cudaFree` 跨所有 stream 同步；`cudaMallocAsync`/`cudaFreeAsync` 排入 stream、不阻塞；忽略 current device，依 pool/stream 決定 device
Memory pool	default/implicit（non-migratable、永遠該 device 可存取、不支援 IPC）vs explicit（`cudaMemPoolCreate`，可 IPC、最大池、CPU NUMA）
Multi-GPU	`cudaMemPoolSetAccess` + `cudaDeviceCanAccessPeer`（不跟隨 peer access，影響 pool 內所有 allocation）
IPC	兩步：先共享 pool（`cudaMemPoolExportToShareableHandle`/`ImportFromShareableHandle`）再共享 allocation（`ExportPointer`/`ImportPointer`）；釋放須 importing 先於 exporting
調校	`cudaMemPoolAttrReleaseThreshold`（`UINT64_MAX` 停用自動縮小）、`cudaMemPoolTrimTo`；reuse policy：`FollowEventDependencies`/`AllowOpportunistic`/`AllowInternalDependencies`；查詢 `cudaDevAttrMemoryPoolsSupported`

Cooperative Groups

項目	速查
handle accessor	`thread_rank`（1D 序號）/ `num_threads` / `thread_index`（3D）/ `dim_threads`
implicit groups	`this_thread_block()` / `this_grid()` / `coalesced_threads()` / `this_cluster()`（CC 9.0+）
partition（collective）	`tiled_partition`（固定大小 1D）/ `labeled_partition` / `binary_partition`；放在非全員到達分支 → deadlock/corruption
同步	`sync()` == `__syncthreads()`；`barrier_arrive`/`barrier_wait`（`arrival_token`）；grid sync 須 `cudaLaunchCooperativeKernel`（CC 6.0+；查 `cudaDevAttrCooperativeLaunch`）；CUDA 13 移除 multi-device sync
collective ops	`reduce`（plus/less/greater/bit_and/or/xor，HW 加速 CC 8.0+ 且僅 4B）、`inclusive_scan`/`exclusive_scan`、`invoke_one`/`invoke_one_broadcast`
memcpy_async	global→shared prefetch + `wait`；雙方 ≥4B 對齊（最佳 16B）才非同步

Programmatic Dependent Launch（PDL）

項目	速查
機制	同一 stream，secondary 在 primary 完成前提早 launch，重疊 preamble；CC 9.0+ 才有 overlapping execution
device 函式	primary：`cudaTriggerProgrammaticLaunchCompletion`（未呼叫則 primary 全 block 退出後隱式觸發）；secondary：`cudaGridDependencySynchronize`（等資料 flush 到 global）
launch	secondary 用 `cudaLaunchKernelEx` + `cudaLaunchAttributeProgrammaticStreamSerialization`（`=1`）
安全/Graph	並行是 opportunistic、不保證（依賴並行可 deadlock）；graph edge `cudaGraphDependencyTypeProgrammatic`，port `...PortProgrammatic`/`...PortLaunchCompletion`（`triggerAtBlockStart` 0→Programmatic、1→LaunchCompletion）

→ Stream-Ordered Memory Allocator、Cooperative Groups 深入、Programmatic Dependent Launch 深入

平台功能（Green Contexts / Lazy Loading / Error Log）

Green Contexts

項目	速查
本質	輕量 context，建立時綁定特定 GPU 資源（SMs + work queues）；只改 host 端、不改 kernel；CUDA 13.1 起以 execution context (EC) 暴露於 runtime
device resource	`cudaDevResourceTypeSm` / `WorkqueueConfig` / `Workqueue`；`cudaDevResource`/`cudaDevResourceDesc_t`（stream 只能關聯 SM 型）
建立四步驟	`cudaDeviceGetDevResource` → `cudaDevSmResourceSplit(ByCount)` → `cudaDevResourceGenerateDesc` → `cudaGreenCtxCreate`（flags=0）；launch 用 `cudaExecutionCtxStreamCreate`
數值/旗標	`coscheduledSmCount`（cluster，CC 9.0+）、`preferredCoscheduledSmCount`（CC 10.0+）、`cudaDevSmResourceGroupBackfill`；`CUDA_DEVICE_MAX_CONNECTIONS` 影響 WQ 數
保證	不保證真正並行，只移除阻礙並行的因素；vs MIG/MPS：GC 固定特定 N 顆 SM、最輕量、允許 SM oversubscription（MPS 靜態不允許）

Lazy Loading 與 Error Log

項目	速查
Lazy loading 預設	CUDA 12.3 起所有平台預設啟用（11.7 引入、12.2 Linux）；需 runtime ≥11.7 AND driver ≥515；不需 compiler；含 managed variable 的 module 仍 eager
控制/查詢	`CUDA_MODULE_LOADING=LAZY`/`EAGER`；`cuModuleGetLoadingMode`（先 `cuInit`）；強制載入 `cuModuleGetFunction`/`cudaFuncGetAttributes`（`cuModuleLoad` 不保證）
三大陷阱	concurrent kernel 序列化 deadlock / 開機吃滿 VRAM 配置失敗 / benchmark 被初始化污染 → 解法：preload kernel / EAGER / `cudaMallocAsync` / warmup
Error Log	設 `CUDA_LOG_FILE`（stdout/stderr/路徑）；格式 `[Time][TID][Source][Severity][API] Message`；僅 CUDA Driver；`cuLogsRegisterCallback`、`cuLogsCurrent`/`cuLogsDumpToFile`/`cuLogsDumpToMemory`（buffer 上限 100 筆、dump memory 上限 25600 bytes、flags 須 0）

→ Green Contexts、Lazy Loading 與 Error Log

非同步機制（Async Barriers / Pipelines / LDGSTS / TMA / STAS）

Asynchronous Barriers

項目	速查
API	`cuda::barrier`；`init(&bar, count)` 設 expected arrival count；split arrive/wait：`arrive()` 不阻塞回 `arrival_token`、`wait(move(token))` 才阻塞
phase	countdown 歸零自動原子 reset 進下一 phase；explicit phase tracking 用 parity（even=0/odd=1）`mbarrier_try_wait_parity`（限 shared、block/cluster scope）
其他	warp 收斂更新 1 次/發散 32 次（先 `__syncwarp`）；`arrive_and_drop`；completion function；transaction barrier CC 9.0+ shared block/cluster，transaction count（bytes）`barrier_arrive_tx`/`barrier_expect_tx`；雙緩衝 producer-consumer = 每 buffer 2 barrier（共 4）

Pipelines

項目	速查
物件	`cuda::pipeline`；非 thread scope 需 `pipeline_shared_state<scope, count>`（count = 緩衝深度）；unified（thread 同時 producer+consumer）vs partitioned（固定角色，thread-local 不能 partition）
流程	submit：`producer_acquire`→`memcpy_async`→`producer_commit`（資源用盡 acquire 阻塞）；consume：`consumer_wait`（等 tail/最舊）→`consumer_release`；`wait_prior<N>`
primitives	`__pipeline_memcpy_async`/`__pipeline_commit`/`__pipeline_wait_prior(N)`；diverge 致 over-wait（先 `__syncwarp`）；提前離開須 `quit()`

Async Copies：LDGSTS / TMA / STAS

機制	速查
LDGSTS	CC 8.0+；僅 global → shared，繞過 register；4/8B = L1 ACCESS、16B = L1 BYPASS；對齊 4/8/16B（最佳 128B）；每 thread 只等自己的複製 → 共享須 `__syncthreads()`；僅 C primitives 保證用 LDGSTS
TMA	CC 9.0 (Hopper)+；bulk-async（1D，免 tensor map）/ bulk-tensor（≤5D，需 tensor map）；讀用 shared barrier transaction count、寫用 bulk async-group（commit/wait）；`CUtensorMap` 由 host `cuTensorMapEncodeTiled`（最快變動維排第一）；傳遞首選 `const __grid_constant__`；單一 thread 發起（`is_elected`/`invoke_one`）；swizzle none/32B/64B/128B；多維 shared 對齊 128B（swizzle 128B→1024B）；device 編碼 `tensormap_replace`（sm_90a* 專屬，rank 零基）
STAS	CC 9.0+（cluster）；唯一方向 register → distributed shared memory；4/8/16B（依大小對齊）；僅低階 `cuda::ptx::st_async`（無高階包裝）；完成靠 mbarrier；遠端 barrier 用 `space_cluster`、本地用 `space_shared`；`cluster.map_shared_rank`

→ Asynchronous Barriers 深入、Pipelines 深入、非同步複製：LDGSTS、非同步複製：TMA、非同步複製：STAS

進階記憶體與排程（Work Stealing / L2 Cache / Sync Domains / IPC / VMM / EGM）

Work Stealing 與 Cluster Launch Control

項目	速查
引入	Cluster Launch Control 為 Blackwell（CC 10.0）；結合 Fixed Work（load balance/preemption）+ Fixed Number（reduced overhead）
機制	block 取消「尚未開始」的 block → 竊取其 index 做事；失敗原因 = 無剩餘 index 或高優先權 kernel；非同步 + mbarrier，async proxy fence；`expect_tx` 用 `sizeof(uint4)`
API/規則	libcu++ `clusterlaunchcontrol_try_cancel`（用 `invoke_one`）；觀察失敗後再 request = UB；cluster 變體 `try_cancel_multicast`（`scope_cluster`、`cluster_group::sync()`、加 `block_index().x`）

L2 Cache Control

項目	速查
需求/API	CC 8.0+；CUDA runtime API（11.0）/ `cuda::annotated_ptr`（11.5）
set-aside	`cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize, size)`，上限 `cudaDeviceProp::persistingL2CacheMaxSize`；MIG 停用、MPS 改用 `CUDA_DEVICE_DEFAULT_PERSISTING_L2_CACHE_PERCENTAGE_LIMIT`
access policy window	`base_ptr`/`num_bytes`(< `accessPolicyMaxWindowSize`)/`hitRatio`/`hitProp`/`missProp`；掛 stream（`cudaStreamAttributeAccessPolicyWindow`）或 graph node；三屬性 Streaming/Persisting/Normal；`hitRatio<1.0` 避免 thrashing
reset	Normal 屬性窗口 / `cudaCtxResetPersistingL2Cache()` / 自動（不建議依賴）

Memory Synchronization Domains

項目	速查
需求	CC 9.0 (Hopper) + CUDA 12.0；Hopper 4 domains，pre-9.0 回報 1（可攜）
解決	fence interference（cumulativity 致保守等待）；每 launch 取 domain ID，fence 只 order 同 domain writes；跨 domain 需 system-scope、同 domain device-scope 即可
API	`cudaLaunchAttributeMemSyncDomain`（logical `Default`/`Remote`）、`cudaLaunchAttributeMemSyncDomainMap`；`cudaDevAttrMemSyncDomainCount`；預設 default→0、remote→1；NCCL 2.16+ 自動標 remote

Interprocess Communication（IPC）

項目	速查
問題	device pointer/event handle 僅在建立它的 process 有效；交換 handle 而非 pointer
Legacy IPC	`cudaIpcGetMemHandle`/`cudaIpcOpenMemHandle`；僅 Linux、不支援 `cudaMallocManaged`、收發 driver/runtime 須一致；子分配風險 → 建議 2 MiB 對齊；Tegra 僅 event-sharing
VMM IPC / 多節點	VMM API 逐 allocation 控制、跨 OS（需 Driver API）；多節點 NVLink 用 fabric handle

Virtual Memory Management（VMM）

項目	速查
定位	Driver API、需 UVA；把 VA 保留與實體配置分離（vs `cudaEnablePeerAccess` 映射所有配置）
五步驟	`cuMemCreate`（回 `CUmemGenericAllocationHandle`，非指標、尚不可存取，size 對齊 granularity）→ export/import → `cuMemAddressReserve` + `cuMemMap` → `cuMemSetAccess`（不設會 crash）→ 釋放 `cuMemUnmap`→`cuMemRelease`→`cuMemAddressFree`（嚴格此序）
handle	POSIX fd/Win32（單節點）vs `CU_MEM_HANDLE_TYPE_FABRIC`（12.4+、單/多節點、需 IMEX）
進階	multicast：`cuMulticastCreate`/`AddDevice`/`BindMem`，`multimem` PTX、NVLink SHARP（CC 9.0+）；compressible memory；virtual aliasing 須 `fence.proxy.alias`

Extended GPU Memory（EGM）

項目	速查
本質	NVLink-C2C，integrated CPU-GPU（Arm）系統；GPU 以 NVLink 速度存取整個系統記憶體；遠端走 GPU NVLink，可跨 NVSwitch
位置/配置	用 OS 的 numaID（≠ device ordinal），`cuDeviceGetAttribute`+`CU_DEVICE_ATTRIBUTE_HOST_NUMA_ID`；allocator：`cuMemCreate`（`CU_MEM_LOCATION_TYPE_HOST_NUMA`）/ `cudaMemPoolCreate`（`cudaMemLocationTypeHostNuma`）
多節點/陷阱	跨節點用 `CU_MEM_HANDLE_TYPE_FABRIC` + export/import；2MB pages；限制裝置用 `CUDA_VISIBLE_DEVICES`（勿用 cgroups，會斷路由）

→ Work Stealing 與 Cluster Launch Control、L2 Cache Control、Memory Synchronization Domains、Interprocess Communication、Virtual Memory Management、Extended GPU Memory

進階執行與互通（Dynamic Parallelism / Graphics / External Interop / Driver Entry Point）

CUDA Dynamic Parallelism

項目	速查
版本	CDP2 = CUDA 12.0+ 預設、CC 9.0+ 唯一；CDP1 為 legacy（`-DCUDA_FORCE_CDP1_IF_SUPPORTED`，將移除）；混用 → `cudaErrorCdpVersionMismatch`
機制	device code 以 `<<<>>>` 啟動 child grid；parent/child 正確巢狀、implicit sync；CDP2 無 `cudaDeviceSynchronize` → `cudaStreamTailLaunch`（等 child）/ `cudaStreamFireAndForget`
記憶體	global/mapped/texture(唯讀) 可傳指標；local/shared 不可（`__isGlobal()` 判斷）；weak consistency（唯一一致點 = child 被 invoke 當下）
Stream/事件	named stream 全 grid 共用、NULL stream 僅 block 內；無新並行保證；events 僅 `cudaStreamWaitEvent`（須 `cudaEventDisableTiming`）
編譯/PTX	`nvcc -rdc=true ... -lcudadevrt`；`cudaLimitDevRuntimePendingLaunchCount`；PTX `cudaLaunchDevice`+`cudaGetParameterBuffer`（buffer 64B 對齊、不可重排、上限 4KB）

Graphics Interoperability

項目	速查
適用	OpenGL 與 Direct3D 9/10/11（不含 D3D12）；`cudaGraphicsResource`；流程 register → map → 取址 → kernel → unmap → unregister（register 昂貴、per-context）
取址	buffer → `cudaGraphicsResourceGetMappedPointer`（device pointer）；texture/array → `cudaGraphicsSubResourceGetMappedArray`（CUDA array）
註冊	GL：`cudaGraphicsGLRegisterBuffer`（pointer）/ `cudaGraphicsGLRegisterImage`（array，寫入需 `...RegisterFlagsSurfaceLoadStore`）；D3D：`cudaGraphicsD3D11RegisterResource`（僅 `D3D_DRIVER_TYPE_HARDWARE`）
SLI	僅 explicit SLI；配置會放大其他 GPU 記憶體；資源綁註冊裝置；`cudaGLGetDevices`/`cudaD3D11GetDevices`

External Resource Interoperability

項目	速查
適用	Direct3D 11/12、Vulkan、NVSCI；OS handle（fd / NT / D3DKMT / NvSciObj）零複製共享
記憶體	`cudaImportExternalMemory` → `cudaExternalMemoryGetMappedBuffer`/`...MipmappedArray`；`cudaDestroyExternalMemory`（不釋放映射，另用 `cudaFree`/`cudaFreeMipmappedArray`）
同步物件	`cudaImportExternalSemaphore` → `cudaSignal/WaitExternalSemaphoresAsync`；`cudaDestroyExternalSemaphore`
配對/規則	Vulkan UUID / D3D12 LUID / NVSCI GPU id；Linux fd 匯入後 CUDA 接管、Windows NT handle 自行 `CloseHandle`；D3D12 須 `cudaExternalMemoryDedicated`；binary（值恆 0）vs timeline semaphore；NVSCI mipmap level 須為 1

Driver Entry Point Access

項目	速查
用途	CUDA 11.3 起取 driver function 指標（類比 `dlsym`/`GetProcAddress`）；舊 toolkit + 新 driver 存取新功能
API	Driver：`cuGetProcAddress(name, &pfn, version, flags, &status)`；Runtime：`cudaGetDriverEntryPointByVersion`；typedef header `cudaTypedefs.h`（`PFN_xxx_vNNNNN`，NNNNN = 引入該符號的 CUDA 版本）
版本規則	version 須精確對應 typedef（傳更高可能換回更新符號 = UB；過低 → `CUDA_ERROR_NOT_FOUND`）；勿用 `CUDA_VERSION`/`cuDriverGetVersion` 當引數
per-thread stream	符號帶 `_ptsz`/`_ptds`；`CU_GET_PROC_ADDRESS_PER_THREAD_DEFAULT_STREAM` / `..._LEGACY_STREAM`（Runtime `cudaEnablePerThreadDefaultStream`/`cudaEnableLegacyStream`）
失敗診斷	`CUresult`（API/usage 錯誤）vs `CUdriverProcAddressQueryResult`：`VERSION_NOT_SUFFICIENT`（升 cudaVersion 即可）/ `SYMBOL_NOT_FOUND`（driver 太舊或拼錯）

→ CUDA Dynamic Parallelism、Graphics Interoperability、External Resource Interop、Driver Entry Point Access

必記重點 / 規則

2003 = pipeline 部分可程式化（僅繪圖）；2006 = CUDA 讓任意運算用 GPU，勿混淆。 → GPU 運算基礎
GPU 用 throughput 換 latency：不讓單一 thread 更快，靠海量並行堆總效能；序列工作 CPU 仍較佳。 → GPU 運算基礎
應用永遠從 host (CPU) 開始執行；host/device 是邏輯角色，SoC 上可能同封裝。 → 執行模型與 SIMT
warp = 32 threads（lane 0–31）；block thread 數應為 32 的倍數，否則最後 warp 有閒置 lane。 → 執行模型與 SIMT
block 間無排程順序保證、不可有資料相依；跨 block 同步須用 cluster（需 CC 9.0+）。 → 執行模型與 SIMT
block = 執行單位、tile = 資料單位；array 可變、tile 不可變且各維為 2 的次方且編譯期已知、不可當 kernel 參數。 → Tile 程式設計
load 越界補零、store 越界丟棄（不對稱）；不同 shape 自動 broadcast。 → Tile 程式設計
register 是 per-thread、shared memory 是 per-block；block 所需 register 超過 register file 則 kernel 無法啟動。 → GPU 記憶體階層
L1 per-SM、L2 全 GPU 共享；mapped memory 走 PCIe/NVLINK，非高效替代。 → GPU 記憶體階層
CC X.Y 直接對應 SM 版本：CC 12.0 → sm_120；compute_XY=PTX、sm_XY=cubin。 → CUDA 平台
cubin：同 major 內 minor >= 目標可載入、跨 major 不相容；PTX 可 JIT 到相同或更高 CC（不可降階），由 device driver 執行、快取於 compute cache。 → CUDA 平台