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

第五章：Technical Appendices

計算能力與環境變數

Compute Capability (CC)

項目	速查
CC 定義/取得	`major.minor` 決定特性與技術規格；`nvidia-smi --query-gpu=name,compute_cap`、`cudaDeviceGetAttribute`、`cuDeviceGetAttribute`、`nvmlDeviceGetCudaComputeCapability`（NVML 需 `-lnvidia-ml`）
三層 feature set	architecture-specific（後綴 `a`，CC 9.0+，僅 exact CC）⊇ family-specific（後綴 `f`，CC 10.0+，同 family）⊇ baseline（`compute_NN`，CC NN.x 及之後）
`compute_100`/`100f`/`100a`	全 10.0+ / 同 family（10.0、10.3）/ 僅 10.0；`compute_120f` 相容 12.0、12.1
Tensor Core 特性	經 inline PTX 提供，官方建議改用 CUDA-X 函式庫（cuDNN/cuBLAS/cuFFT）或 CUTLASS
不變上限	warp size 32、threads/block 1024、registers/thread 255、block z-dim 64、grid x-dim 2³¹−1、grid y/z 65535、resident grids/device 128
記憶體不變上限	32-bit registers/SM 64K、registers/block 64K、shared banks 32、local/thread 512KB、constant 64KB、constant cache/SM 8KB
shared/SM 例	7.5=64KB、8.0/8.7=164KB、8.6/8.9/12.x=100KB、9.0/10.x/11.0=228KB
>48KB shared/block	須改用 dynamic shared memory 並明確 opt-in
8.x baseline 例	Warp Reduce、Bfloat16、硬體 `memcpy_async`、Async Barriers、L2 residency、DPX（多指令模擬）
9.0 baseline 例	Thread Block Cluster、Distributed Shared Memory、TMA、128-bit 整數 atomics、native DPX、`atomicAdd` float2/float4
位元組序	所有 NVIDIA GPU 為 little-endian
Tensor Core 輸入型別	7.5=FP16/INT8/INT4；8.0 加 FP64/TF32/BF16；8.9 加 FP8；10.0+ 加 FP6/FP4

CUDA 環境變數

項目	速查
`CUDA_VISIBLE_DEVICES`	控制可見 GPU 與列舉順序（決定 ordinal）；未設=全可見、空字串=全不可見；三形式 integer index / GPU UUID（可縮寫前綴）/ MIG（僅單一 instance）；遇無效 index 後面截斷不可見
`CUDA_DEVICE_ORDER`	`FASTEST_FIRST`（預設）/ `PCI_BUS_ID`
`CUDA_MANAGED_FORCE_DEVICE_ALLOC`	非零=強制 Unified Memory 用 device memory；要求所有 device P2P 相容否則 `cudaErrorInvalidDevice`
`CUDA_CACHE_DISABLE`/`_PATH`/`_MAXSIZE`	JIT 快取開關/路徑/大小；預設桌機/伺服器 1GiB、嵌入式 256MiB、上限 4GiB；Linux 路徑 `~/.nv/ComputeCache`
JIT override	`CUDA_FORCE_PTX_JIT` 覆蓋 `CUDA_FORCE_JIT`；`CUDA_DISABLE_PTX_JIT` 覆蓋 `CUDA_DISABLE_JIT`；`CUDA_FORCE_PRELOAD_LIBRARIES=1` 避免多執行緒 deadlock
`CUDA_LAUNCH_BLOCKING`	=1 關非同步，使錯誤對齊觸發的 API 呼叫（除錯用、會變慢）
`CUDA_DEVICE_MAX_CONNECTIONS`	work queue 數 1–32，預設 8；不足→false dependency 序列化
`CUDA_DEVICE_MAX_COPY_CONNECTIONS`	僅 CC 8.0+；兩者都設時覆蓋前者的 copy connections
`CUDA_DEVICE_WAITS_ON_EXCEPTION`	=1 device exception 時暫停，可 attach `cuda-gdb` 檢視 live state
`CUDA_SCALE_LAUNCH_QUEUES`	launch queue（command buffer）大小縮放；合法值 `0.25x`/`0.5x`/`2x`/`4x`，其餘一律當 `1x`
`CUDA_GRAPHS_USE_NODE_PRIORITY`	graph 執行優先序：=0 繼承 stream（預設）/ =1 用 per-node 優先序當排程提示（覆蓋 `cudaGraphInstantiateFlagUseNodePriority`）
`CUDA_DEVICE_DEFAULT_PERSISTING_L2_CACHE_PERCENTAGE_LIMIT`	persisting L2 保留比例 0–100（預設 0）；僅 CC 8.0+ 且用 MPS，須在 `nvidia-cuda-mps-control -d`（啟 MPS）前設
`CUDA_DISABLE_PERF_BOOST`	=1 不提升 device performance state、改 heuristic 隱式選 pstate；僅 Linux
`CUDA_MODULE_LOADING`	`LAZY`（預設，降 startup/記憶體、首呼叫載入）/ `EAGER`（全載、launch 開銷可預測）；`CUDA_MODULE_DATA_LOADING` 為資料層互補設定
`CUDA_BINARY_LOADER_THREAD_COUNT`	載入 device binary 用的 CPU thread 數；0=採預設 1 個 thread
`CUDA_LOG_FILE`	描述性錯誤輸出至 `stdout`/`stderr`/檔案路徑
`CUDA_AUTO_BOOST`	已棄用，改 `nvidia-smi --applications-clocks` / NVML

→ 計算能力、CUDA 環境變數

C++ 語言支援（特性、標準庫、限制）

C++ 標準特性支援（nvcc --std）

項目	速查
標準/起始 Toolkit	C++11=7.x、C++14=9.x、C++17=11.x、C++20=12.x；`-std=c++NN` 開啟該版全部特性並驅動 host preprocessor/compiler/linker
C++11 device 不支援	整組 Concurrency（`std::atomic`、memory model、bidirectional fences、`thread_local`、propagating exceptions）+ GC 最小支援(N2670) + extended integral types(N1988)
C++14 不支援	僅 Clarifying memory allocation(N3664)、Sized deallocation(N3778)
C++17	列表內全支援（structured bindings、`constexpr if`、constexpr lambda、fold expressions、inline variables）
C++20 支援/不支援	支援 Concepts、`operator<=>`、`consteval`、`constinit`、`char8_t`；不支援 Modules 與 Coroutines
C++20 host compiler	GCC≥10、Clang≥10、MSVS≥2022、nvc++≥20.7
C++23	支援 `Deducing this`、`if consteval`、`[[assume]]`；5 項 × + 3 項 N/A；host compiler GCC≥14、Clang≥18、MSVC 不支援、nvc++≥24.3
libcu++	CUDA 版 C++ STL（`cuda::std::`，`<cuda/std/...>`），host+device；C++20/23/26→C++17 backport；extended types `__int128`/`__half`/`__nv_bfloat16`/`__float128`；隨 Toolkit 或 CCCL

C standard library、Lambda 與 nvstd::function

項目	速查
device `printf()`	每 thread 各自執行；回傳解析引數數（非字元數，無引數 0、NULL −1、內部錯誤 −2）；上限 32 引數；host 端格式化；1MB 環狀 buffer（`cudaLimitPrintfFifoSize`），程式結束不 flush
device `malloc/free`	device heap，回 16-byte 對齊；預設 heap 8MB（`cudaLimitMallocHeapSize`，須配置前設、module 載入後不可變）；與 `cudaMalloc`/`cudaFree` 互不相通
`clock()`/`clock64()`	per-multiprocessor 計數器；因 time-slice 量到值大於實際週期
`alloca()`	caller stack frame 配置、回傳自動釋放；device 16-byte 對齊
Lambda 執行空間	由最內層 enclosing function 決定；無 scope→`__host__`；傳給 `__global__` 須 `__device__` 或 `__host__ __device__`；global/namespace scope lambda 不可當引數
Extended lambda	定義在 `__host__`/`__host__ __device__` 函式內、帶執行空間標註；需 `--extended-lambda`（巨集 `__CUDACC_EXTENDED_LAMBDA__`）；只能 by value 捕捉、陣列維度≤7、host-device 不可 init-capture；存在/順序不可依賴 `__CUDA_ARCH__`
`*this` 捕捉(C++17)	複製物件本身而非 `this` 指標，避免 GPU 存取 host `this->member` 的 runtime error
`nvstd::function`	`<nvfunctional>`；host/device 皆可用但不可跨邊界傳遞/初始化；`operator()` 僅 `__device__`

C/C++ device 語言限制

項目	速查
device 不支援	RTTI(`typeid`/`dynamic_cast`)、例外(`try/catch/throw`)、`long double`、trigraphs；自訂 `operator new/delete` 取代內建為 UB
保留 namespace	`cuda::`/`nv::`/`cooperative_groups::` 加定義是 UB（唯一例外：`cuda` 巢狀於非保留 namespace）
跨空間存取	host 解 device 指標、device 解 host 指標皆 UB；`cudaGetSymbolAddress` 取得位址只能 host 用
recursion	`__global__`/`__tile_global__`/`__tile__` 不支援；`__device__`、`__host__ __device__` 可
`__global__` 參數	經 constant memory，總大小上限 32,764 bytes；不可 variadic/`std::initializer_list`；polymorphic class 參數為 UB
host→device 傳參	raw memcpy 跳過 copy constructor；destructor 可能在 kernel 完成前執行
const(host)變數值入 device	須 constant expression 初始化、非 volatile、且為 built-in 整數/浮點（MSVC 浮點除外）；建議改 `constexpr`
`volatile`	非 atomic、不保證順序/存取次數；不適執行緒同步（改 atomic）或 MMIO（改 PTX MMIO）
static 變數	`__host__ __device__` 函式各空間一份；不允許 dynamic initialization
tile code 不支援	`switch`/`goto`/迴圈內 `return`、virtual 呼叫、function pointer、bitfield、128-bit 型別、>16MB 型別
`--expt-relaxed-constexpr`	放寬 constexpr 跨 execution space 呼叫，定義 `__CUDACC_RELAXED_CONSTEXPR__`

→ C++ 語言特性支援、C 標準函式庫、Lambda 與多型包裝器、C/C++ 語言限制

語言擴充：限定詞、同步、原子、Warp

函式與變數限定詞

項目	速查
執行空間	`__host__`(預設)/`__device__`/`__global__`(host+SIMT 啟動、回 void、非 member、不可遞迴、需 `<<<>>>`、非同步)/`__tile__`/`__tile_global__`；多空間用 `__CUDA_ARCH__` 分流（host 端未定義）
記憶體空間	`__device__`(全域、每 device 一份)/`__constant__`(device 唯讀、host 改)/`__shared__`(block、static 或 `extern` 動態)/`__managed__`(host/device 自動、位址非常數運算式)
host 存取 symbol	`cudaGetSymbolAddress/Size`、`cudaMemcpyTo/FromSymbol`
`__restrict__`	承諾不別名放開優化；增暫存器壓力；const `__global__` 指標→唯讀快取載入(`ld.global.nc`/`__ldg`)
`__grid_constant__`	`__global__` const 非 reference 參數共用單一位址、唯讀、kernel 生命週期
inlining	`__noinline__`/`__forceinline__`/`__inline_hint__`(LTO) 三者互斥
`__launch_bounds__`	`maxThreadsPerBlock`(必要、`.maxntid`)/`minBlocksPerMultiprocessor`/`maxBlocksPerCluster`；引導暫存器與 occupancy；超過則 launch 失敗
`__maxnreg__`	單 thread 最大暫存器(`.maxnreg`)；不可與 `__launch_bounds__` 並用；含它的 kernel 忽略 `--maxrregcount`
`__cluster_dims__`	CC 9.0+ 編譯期 cluster 維度；不可用於 `__tile_global__`
內建變數	`gridDim`/`blockDim`=`dim3`，`blockIdx`/`threadIdx`=`uint3`，`warpSize`=`int`；`dim3` 分量預設 1
向量型別	`make_<type>()`；`.x/.y/.z/.w`；`int4`/`float4`/`double2` size16 align16；`int3`/`float3` size12 align4；CUDA 13 deprecate `long4`/`ulong4`/`longlong4`/`ulonglong4`/`double4`
execution config	`<<<grid_dim, block_dim, dynamic_smem_bytes, stream>>>`；`__tile_global__` 的 block_dim 須為 1；config 引數先於函式引數求值

同步原語與 Memory Fence

項目	速查
記憶體模型	weakly ordered；無 fence/sync 對同址讀寫=UB
`__syncthreads()`	等 block 內所有未退出 thread；strongly happens-before；條件呼叫須 block uniform 否則 hang
`__syncthreads_count/and/or`	同步+predicate 規約：非零數量 / 全為真 / 至少一真
`__syncwarp(mask)`	同 warp mask 指名 thread 同步，預設 mask `0xFFFFFFFF`
Fence 三層	`__threadfence_block`(block)/`__threadfence`(device)/`__threadfence_system`(device+host+所有 peer device)；只排序不保證可見性（可見性靠 volatile/atomic）
跨 block 規約	最後完成的 block 由 `atomicInc(&count, gridDim.x)` 取得舊值 `gridDim.x − 1` 判定
libcu++ 推薦	fence 用 `cuda::atomic_thread_fence`；原子用 Extended `cuda::atomic`/`cuda::atomic_ref`

原子函式

項目	速查
五種提供方式	Extendedatomic，無 scope)/Built-in(`__nv_atomic_*`，CUDA 12.8+、device only)/Tile/Legacy(`atomic<Op>`，僅 relaxed、不引入 fence
Legacy scope 後綴	無=device、`_block`=block、`_system`=system
`atomicInc(addr,val)`	`old >= val ? 0 : old+1`（環繞）
`atomicCAS`	`old == compare ? val : old`，回傳 old；萬用基底（任何原子運算可用 CAS 迴圈實作）
型別 CC 門檻	`atomicAdd` 的 `__nv_bfloat16/162` 需 CC 8.x+；`float2/float4` 需 CC 9.x+ 且僅 global；`atomicExch`/`atomicCAS` 128-bit 模板需 CC 9.x+（align≥16、trivially copyable）
`atomicExch` 型別	`int`/`unsigned`/`unsigned long long`/`float`（無 `double`，`double` 只 `atomicAdd`）；`unsigned short` 僅 `atomicCAS`
向量原子性	逐元素原子，整個向量非單次原子
浮點 CAS 迴圈	須用整數位元比較，避免 `NaN != NaN` 無法收斂 hang
`__nv_atomic_*` 限制	device only、不可 local memory、不可取位址；order/scope 須整數字面量；`weak` 參數被忽略（取較強者）

Warp 函式與裝置函式

項目	速查
`__activemask()`	active lane 瞬間快照；不等於走某分支的 lane；不可拿去當 `*_sync` 的 mask
Warp vote	`__all_sync`(全非零)/`__any_sync`(至少一)/`__ballot_sync`(第 N bit=predicate 非零且 active)
Warp match	`__match_any_sync`(回同值 lane 遮罩)/`__match_all_sync`(全同回 mask 並設 pred)
Warp reduce	`__reduce_<op>_sync`，CC 8.x+；add/min/max + and/or/xor
Warp shuffle	`__shfl_sync`/`_up`/`_down`/`_xor`；不經 shared memory；`width` 須 2 的次方且 ∈[1,32]；讀 inactive thread=未定義
mask 鐵則	呼叫者 bit=1、非呼叫者 bit=0、退出忽略；參與者用相同 mask（disjoint 例外）；warp `_sync` 不提供 memory ordering*
`__CUDA_ARCH__`	僅 device code；值=`compute_<version>*10`（compute_80→800）
arch-specific 巨集	`__CUDA_ARCH_SPECIFIC__`/`__CUDA_ARCH_FAMILY_SPECIFIC__` 對應 nvcc `compute_<v>a`/`compute_<v>f`，僅 device code；呼應筆記01 三層 feature set（如 `compute_100f/sm_103f`→`_SPECIFIC__` 未定義、`_FAMILY_SPECIFIC__`=1000）
Feature Testing 巨集	nvcc 支援該特性時定義：`__CUDACC_DEVICE_ATOMIC_BUILTINS__`(device atomic builtins)/`__NVCC_DIAG_PRAGMA_SUPPORT__`(診斷 pragma)/`__CUDACC_EXTENDED_LAMBDA__`(extended lambda)/`__CUDACC_RELAXED_CONSTEXPR__`(relaxed constexpr)
`__nv_pure__`	標記 pure 函式（無副作用、可讀不改全域）；host+device；轉 GNU `pure`/MSVC `noalias`
位址空間	`__isGlobal`/`__isShared`/`__isConstant`/... 判斷；`__cvta_generic_to_`(回 `size_t`)/`__cvta__to_generic`(回 `void*`)
低階載入	`__ldg`(read-only L1/Tex)；`__ldcg`/`__ldca`/...、`__stwb`/... 依 PTX 快取運算子
`__trap()`/`__nanosleep`	`__trap` 中止 kernel、損毀 context；`__nanosleep(ns)` 約睡 ns、上限約 1ms
DPX	3 路 min/max + fused add-min/max + ReLU（夾到 0）；Smith-Waterman/Needleman-Wunsch/Floyd-Warshall 等動態規劃

→ 函式與變數限定詞、同步原語與原子函式、Warp 函式、CUDA 巨集與裝置函式

編譯器提示、診斷與 Warp Matrix

最佳化提示與診斷

項目	速查
`#pragma unroll`	緊接迴圈正前方；無引數=trip count 常數時完全展開；0/1=不展開；非正或 >`INT_MAX`=忽略並警告
`__builtin_assume(pred)`	假設為真；執行期為假=UB、有副作用=unspecified（`__assume` 僅 MSVC）
`__builtin_expect(in,exp)`	分支預測提示、回傳 `in`；類似 `[[likely]]`/`[[unlikely]]`
`__builtin_assume_aligned`	`align` 須 2 的次方且整數字面值；三參數版假設 `(char*)ptr-offset` 對齊
`__builtin_unreachable()`	告知控制流不會到達；到達=UB
`#pragma nv_abi`	separate compilation 保留暫存器；call site 只對間接呼叫有效（僅自由函式）
`#pragma nv_mma_throughput`	MMA 專門最佳化；僅 entry-function scope（`__global__`），效果可沿呼叫路徑傳到被呼叫的 `__device__`；實驗性、不保證每個 kernel 受益
`assert()`	條件為 0 中止 kernel、損毀 context（後續回 `cudaErrorAssert`）；`NDEBUG`/`-DNDEBUG` 停用；同步函式才印訊息
`__brkpt()`	任一 thread 呼叫即暫停 kernel
`nv_diag_*` pragma	suppress/warning/error/default/once；只 warning 可被抑制、任何診斷可升 error；只影響 nvcc 前端（巨集 `__NVCC_DIAG_PRAGMA_SUPPORT__`）

Warp Matrix (WMMA)

項目	速查
概念	`nvcuda::wmma`，Tensor Core 做 `D=AB+C`；CC 7.0+*、全 warp 協作；條件式須整 warp 求值相同否則 hang
API	`fragment`/`load_matrix_sync`/`store_matrix_sync`/`fill_fragment`/`mma_sync`
fragment Use	`matrix_a`(m×k)/`matrix_b`(k×n)/`accumulator`(m×n)；a/b 須指定 `row_major`/`col_major`，accumulator 保留 `void`
對齊/步幅	`mptr` 須 256-bit 對齊；`ldm` `__half` 倍數 8、`float` 倍數 4（皆 16 bytes 倍數）；所有參數須全 warp 一致
型別/尺寸	`__half`、`char`/`unsigned char`(→`int`) 皆 16x16x16/32x8x16/8x32x16（`__half`/整數型 CC 7.0+）；`__nv_bfloat16` shape 同 `__half`、`double` 8x8x4、`tf32` 僅 16x16x8（三者皆需 CC 8.0+）
`satf=true`	+Inf→+MAX_NORM、−Inf→−MAX_NORM、NaN→+0
`double` WMMA	8x8x4，`.rn`（round to nearest even）
sub-byte	`u4`/`s4`/`b1`（experimental）；layout 固定 a=`row_major` b=`col_major`；`bmmaBitOpXOR` 需 compute_75+、`bmmaBitOpAND` 需 CC 8.0+
Restrictions	fragment 是架構特定不透明結構；跨 link-compatible 架構互傳致錯誤/損毀；應先 `store_matrix_sync` 到記憶體再傳指標

→ 編譯器最佳化提示、診斷與 Warp Matrix

浮點數與數學函式

浮點數基礎與 IEEE-754

項目	速查
host/device 差異來源	rounding、FMA contraction、加乘非結合律、編譯器優化、library 實作
預設 rounding	round-to-nearest-ties-to-even(rn)；四種 rn/rz/ru/rd，非預設用 intrinsic（GPU 無動態切換）
denormal	exponent 全 0 的 subnormal；`-ftz=true` flush-to-zero（含於 `--use_fast_math`）
FMA	`ab+c` 只 round 一次、防 subtractive cancellation；`-fmad=true` 啟用 `xy+z` contract
特殊值	`+0==-0` true；`NaN==NaN` false（`NaN!=NaN` true）；`0/0`/`Inf-Inf`/`Inf/Inf`/`Inf*0`/`sqrt(-1)` 得 NaN
IEEE-754 符合	754-2019；無動態 rounding、無例外偵測、SNaN 不 signaling
浮點除 0	implementation-specific 得 infinity（非 UB）；整數除 0 才 UB
可移植旗標	`-ftz=false`/`-prec-div=true`/`-prec-sqrt=true`，不開 `--use_fast_math`（re-association 與 FMA 仍預設允許）
atomic float add denormal	global 永遠 flush-to-zero；shared 永遠支援 denormal（不論 `-ftz`）
型別/CC	half/float/double/quad 是 IEEE-754，Bfloat16 不是；Bfloat16 需 CC 8.0+，`__float128` 需 CC≥10.0
內建運算子	`+ - * / 1/x` 符合 IEEE-754，rn 下 0 ULP；`-prec-div=false` 後 `/`→2 ULP、`1/x`→1 ULP
device 數學	不設 errno、不回報浮點例外

數學函式庫與 Intrinsic

項目	速查
標準路徑	`cuda::std::`（`<cuda/std/cmath>`），host+device；對映 Math API、有 max ULP 界（測試得出、不保證）
全 0 ULP 類別	Basic Ops、Nearest Integer、FP Manipulation、Classification/Comparison、單雙精度基本 intrinsic
取整	用 `rint`/`rintf`（單指令）勿用 `round`/`roundf`（多指令）；`round` ties away from zero、`rint` ties to even
intrinsic 三特性	`__` 前綴、只 device、映射較少指令（快但不準）
捨入後綴	`_rn`(最近偶數)/`_rz`(趨零)/`_ru`(+∞)/`_rd`(−∞)
基本 intrinsic	`__fadd_`/`__dadd_`/`__fmul_`/`__dmul_` 等 0 ULP；前述加乘不會被合併成 FFMA/DFMA（`+`/`*` 常被合併）
`fdividef` vs `__fdividef`	`fdividef` 0 ULP 等同 `x/y`；intrinsic `__fdividef` 對 \|y\|∈[2⁻¹²⁶,2¹²⁶] 為 2 ULP
Non-Standard 函式族	非標準擴充、逐函式定 host/device（`__half`/`bf16`/`__float128` 版 device-only）：`rsqrt`(1/√x)/`rcbrt`/`norm3d`(√Σx²)/`norm4d`/`rnorm3d`/`rnorm4d`/`sinpi`(sin πx)/`cospi`/`erfinv`/`erfcinv`/`normcdf`/Bessel(`cyl_bessel_i0`/`i1`、`j0`/`j1`/`y0`/`y1`/`jn`)
`norm(dim,p)`/`rnorm(dim,p)`	device-only；用快速演算法無誤差界（round-off 損精度）
型別命名前綴	`__half`/`__nv_bfloat16` 用 `h`（`hexp`/`hsin`/`hsqrt`）；`__float128` 用 `__nv_fp128_`（`__nv_fp128_exp`/`__nv_fp128_fma`）
ULP 例	`sqrtf` 0（`--use_fast_math` 變 1）；`powf` 4/`pow` 2；`tgammaf` 5/`tgamma` 10；`__expf`=2+⌊\|1.173·x\|⌋
half/bf16 N/A	無原生實作者轉 float 計算再轉回（如 error/gamma 類）
`--use_fast_math`	把一組 device Math API（含 `cuda::std`）換成 intrinsic

→ 浮點數運算基礎、數學函式庫與 Intrinsic 函式

device 端 API：記憶體屏障、Pipeline、Cooperative Groups 與 Device Runtime

Memory Barrier 與 Pipeline Primitives

項目	速查
`__mbarrier_*`	`cuda::barrier` 的 C-like primitives（`<cuda_awbarrier_primitives.h>`）；`bar` 須在 `__shared__`；`init`/`arrive`/`test_wait`/`try_wait`；arrive 前須 init 且 pending count≠0
arrive vs arrive_and_drop	後者額外遞減下一 phase 的 expected count
test_wait vs try_wait	test_wait 立即查詢；try_wait 未完成可能掛起到 phase 完成或 `max_sleep_nanosec`
`*_parity` 系列	`__mbarrier_test_wait_parity`/`__mbarrier_try_wait_parity` 以 phase 奇偶判定：`true`=奇數 phase、`false`=偶數 phase
`__mbarrier_maximum_count`	回傳允許的最大 arrival count；`init` 的 `expected_count` 與 `__pipeline_arrive_on` 的 +1 皆不可超過此上限
重用 mbarrier 前	須先 `__mbarrier_inval(bar)`
`__pipeline_memcpy_async`	global→shared 非同步；`size_and_align` 只能 4/8/16、等於兩端對齊；`zfill` bytes 補零
`__pipeline_commit`/`wait_prior(N)`	提交批次 / 提交索引 `{0..L}` 等到 `L-N`（含）
`__pipeline_arrive_on`	arrival count +1，先前 memcpy_async 全完成後 −1（淨效果 0）

Cooperative Groups

項目	速查
三大 implicit group	`this_thread_block()`/`this_cluster()`(CC 9.0+)/`this_grid()`；collective 須全員參與否則 UB
`tiled_partition<Size>`	1D row-major 子群；`Size` 2 的次方且 ≤1024、父群須整除；native 1/2/4/8/16/32
`coalesced_threads()`	機會性取當下 active threads；`meta_group_size`=1、`meta_group_rank`=0
tile warp 函式限制	`shfl_up`/`shfl_down`/`shfl_xor`/`ballot`/`match_` 僅 ≤32*；`shfl` >32 須相同 `src_rank`
partition	`tiled_partition`/`labeled_partition`(CC 7.0+)/`binary_partition`(label 0/1 特例)
`cg::reduce`	CC 8.0+ 且 add/min/max 或 AND/OR/XOR、僅 4-byte 硬體加速；限 `coalesced_group`/`thread_block_tile`
scan	inclusive 含呼叫者自身；exclusive 只含 `thread_rank` 較低者
`cg::memcpy_async`	global→shared 非阻塞（CC 8.0 才非同步）；`shape` 版以 bytes、`dstLayout`/`srcLayout` 版以元素
`_update`/`_async`	用 `cuda::atomic`/`cuda::atomic_ref` 以 relaxed 原子更新
`grid.sync()`	須 `cudaLaunchCooperativeKernel`（非 `<<<>>>`）；CC 6.0+；查 `cudaDevAttrCooperativeLaunch`

CUDA Device Runtime

項目	速查
用途/引入	kernel 內呼叫的 host runtime 子集；最常用於 Dynamic Parallelism(CDP2) 與 Device Graph Launch；原型編譯自動帶入無需 include
資源上限	host `cudaDeviceSetLimit()`、kernel 啟動前設、執行中不可改；pending launch slot 預設 2048、event slot 為其兩倍
device 配置	`cudaMalloc`/`cudaFree` 映射到 device `malloc`/`free`，受 heap 限制；跨環境配置/釋放是錯誤
stream	只能 `cudaStreamCreateWithFlags(cudaStreamNonBlocking)`；handle 為建立它的 grid 私有；無 `cudaStreamSynchronize`/`Query`
event	須 `cudaEventDisableTiming`
同步替代	無 `cudaDeviceSynchronize`(CDP2)；`cudaStreamTailLaunch` 等子 kernel、`cudaStreamFireAndForget` 立即排程不被相依
memcpy 限制	只 async、只 device-to-device、不可傳 local/shared 指標
`__constant__`	device 端唯讀；symbol 可直接 `&` 取址，故不需 `cudaMemcpyToSymbol`
`%smid`/`%warpid`	volatile，不保證 thread/block 生命週期內不變
其他	`cudaGetLastError` 為 per-thread；launch 無錯≠child 成功；無多 GPU（可查任意 device 屬性）；ECC 錯誤在 launch tree 完成後 host 側回報；fire-and-forget/tail launch 需 64-bit 且不可定義 `CUDA_FORCE_CDP1_IF_SUPPORTED`

→ 記憶體屏障、Pipeline 與 Cooperative Groups API、CUDA 裝置端 Runtime

記憶體模型與執行模型

CUDA C++ 記憶體模型（Thread Scope）

項目	速查
動機	同步成本非均勻：距離越遠越貴；`cuda::` 在標準模型加 thread scope
四級 scope	`thread_scope_system`(多 GPU+CPU)→`_device`(同 device 且同 memory sync domain)→`_block`(同 block)→`_thread`(只自己)
`std::`/`cuda::std::`	等同 `cuda::` 型別以 `thread_scope_system` 實例化
atomic 何時 atomic	指定 scope 非 system 時自動成立；system scope 需裝置屬性：system allocated→`pageableMemoryAccess`、managed→`concurrentManagedAccess`、mapped→`hostNativeAtomicSupported`
GPU memory 跨 GPU	`cudaDeviceGetP2PAttribute`+`cudaDevP2PAttrNativeAtomicSupported`=1，或僅單一 GPU 存取
data race	兩並行衝突動作，至少一個在涵蓋對方執行緒的 scope 下非 atomic，且無 happens-before → UB
barrier/latch/semaphore	成員函式（除解構子）並行呼叫不引入 data race
Message Passing 範例	兩端須同用 `thread_scope_device`；store 誤用 block scope→load 在別 block 不被涵蓋→UB

CUDA C++ 執行模型（Forward Progress）

項目	速查
device 提供	parallel forward progress [intro.progress.9]；前提 host 提供 concurrent FP [intro.progress.7]
concurrent vs parallel	concurrent：未終止前一定推進；parallel：執行第一步前不保證開始，執行一步後升級 concurrent
host thread	implementation-defined；通用平台應提供 concurrent FP
推進傳染範圍	某 device thread 推進後：Cooperative Grid→整個 grid；否則→整個 thread-block cluster（蘊含整 block）；其他 cluster 不保證
device 少的兩項	`std::this_thread::yield`(host.2) 與 trivial infinite loop(host.6)；automatic storage 的 volatile/atomic 不算進度；只認 atomic read
CUDA API 保證	每次呼叫最終 return 或確保至少一 device thread 推進（不必相關）；query 不可在無推進時一直回 `cudaErrorNotReady`
單次 vs 反覆 query	單次不保證 device 啟動；spin-loop 內反覆 query（或一定推進的 sync API）才保證最終推進
dependencies	device thread 在所有依賴完成前不得啟動；跨不同 stream 無依賴→可能 starvation；同 stream 順序依賴→終止
測 FP bug 環境變數	`CUDA_DEVICE_MAX_CONNECTIONS=1 CUDA_LAUNCH_BLOCKING=1`（非充分）

→ CUDA C++ 記憶體模型、CUDA C++ 執行模型

必記重點 / 規則

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 平台