第五章練習題 (Practice - Technical Appendices)

作答前可先用這張表喚起記憶；每列是「情境/關鍵字 → 答案」。

baseline（compute_NN，引入後所有後續 CC 皆可用）⊇ 反過來被含括：architecture-specific（後綴 a，CC 9.0 起，僅 exact CC）⊇ family-specific（後綴 f，CC 10.0 起，同 family）⊇ baseline；即 a ⊇ f ⊇ baseline。
compute_100：所有 CC 10.0 及之後；compute_100f：同 family（10.0 與 10.3）；compute_100a：僅 CC 10.0。
不變上限：warp size=32、max threads/block=1024、max 32-bit registers/thread=255、grid x-dim=2³¹−1、y/z-dim=65535、block z-dim=64。NVIDIA GPU 一律 little-endian。

遇到無效 index（-1）會截斷清單，只有排在它「之前」的 device 可見，故只有實體 2 與 1 可見，cudaGetDeviceCount() 回傳 2（device 3 在 -1 之後不可見）。
列舉順序決定 ordinal：ordinal 0 = 列舉最前的實體 device 2，所以 cudaSetDevice(0) 選到實體 GPU 2。
設為空字串時「沒有任何 GPU 可見」（未設定才是全部可見）。

CUDA_LAUNCH_BLOCKING=1 關閉非同步執行，讓 CUDA API 錯誤「對齊到觸發它的那個呼叫」便於除錯；代價是程式變慢，是除錯工具非效能設定。
CUDA_MODULE_LOADING 預設 LAZY：延後到取得 function handle 才載入，降低 startup 時間與 GPU 記憶體佔用；EAGER 在初始化即全載，launch 開銷可預測但 startup/記憶體較高。
*_PTX_JIT 版永遠覆蓋對應的 *_JIT 版：CUDA_FORCE_PTX_JIT > CUDA_FORCE_JIT、CUDA_DISABLE_PTX_JIT > CUDA_DISABLE_JIT。

用 --std=c++03/11/14/17/20/23；傳入會開啟該版「全部語言特性」，並以對應 dialect 一併驅動 host 的 preprocessor/compiler/linker。
C++11 device code 不支援整個「Concurrency 群組」：std::atomic、memory model、bidirectional fences、thread_local、propagating exceptions（另有 GC 最小支援、extended integral types）。C++20 缺席 Modules 與 Coroutines（Concepts/<=>/consteval/constinit/char8_t 皆可用）。
libcu++ 是 CUDA 版 C++ 標準庫（cuda::std::），host/device 皆可用，提供 device 上 std::atomic/memory model 等的替代品與 extended types（__int128/__half/__nv_bfloat16/__float128）。

device printf() 回傳「解析的引數數」（非字元數）：無引數→0、format 為 NULL→-1、內部錯誤→-2；引數上限 32 個（不含 format string），多出的被忽略。
Buffer 預設 1 MB、固定大小的「環狀」緩衝（cudaLimitPrintfFifoSize 可調），滿了覆蓋舊輸出；只在 kernel launch/同步/blocking memcpy/module load-unload/context 銷毀/host callback 前 flush，「程式結束不會」自動 flush。
device malloc() 從 device heap 配置、回傳 16-byte 對齊；heap 預設 8 MB（須在任何配置前用 cudaLimitMallocHeapSize 設）。device heap 與 host runtime API 記憶體互不相通，「不可」用 cudaFree/cudaMemcpy 操作。

由「最內層 enclosing function scope 的執行空間」決定；無 scope 則為 __host__。只有執行空間為 __device__ 或 __host__ __device__ 的 lambda 能當 __global__ 引數，global/namespace scope lambda 不可。
Extended lambda = 定義在 __host__/__host__ __device__ 函式內、帶執行空間標註的 lambda（需 --extended-lambda）；只能 by value 捕捉、陣列維度 ≤ 7、host-device 不可 init-capture、存在/順序不可依賴 __CUDA_ARCH__。C++17 *this 捕捉複製「物件本身」而非 this 指標，避免 GPU 上存取 host this->member 的 runtime error。
nvstd::function（<nvfunctional>）host/device 皆可用，但「不可」跨 host/device 邊界傳遞或初始化；其 operator() 僅 __device__。

不支援：RTTI（typeid/dynamic_cast）、例外（try/catch/throw）、long double、trigraphs，以及 __global__/__tile_global__/__tile__ 的 recursion。__global__ 參數經 constant memory 傳遞，總大小上限 32,764 bytes，且不可 variadic、不可 std::initializer_list。
陷阱一：CUDA Runtime 用 raw memcpy 複製參數到 device，「跳過 user-defined copy constructor」，其副作用不會發生（如指向自身成員的指標在 device 端失效）。
陷阱二：kernel 啟動與 host 非同步，若參數有 non-trivial destructor，它可能在「kernel 完成前」就於 host 端執行，破壞有副作用的邏輯。

__launch_bounds__(maxThreadsPerBlock[, minBlocksPerMultiprocessor]) 告訴編譯器啟動上限，據此推導暫存器上限 L 並壓低暫存器用量，確保指定數量的 block 能常駐 SM；建議至少加單引數版以避免該錯誤。第一個參數 maxThreadsPerBlock 編成 PTX .maxntid。
__launch_bounds__ 與 __maxnreg__（編成 .maxnreg）「不可」同時套在同一 kernel。
用 __grid_constant__：const-qualified 非 reference 的 __global__ 參數，全 thread 共用單一位址、唯讀、具 kernel 生命週期（修改它為 UB）。

__syncthreads() 等到 block 內「所有未退出 thread」都到達同一呼叫點或已退出，並建立 strongly happens-before 的記憶體排序；放在條件式中只有當條件「對整個 block uniform」時才合法，否則可能 hang 或 UB。
__threadfence_block：呼叫 thread 所在 block 內所有 thread；__threadfence：本 device 內任一 thread；__threadfence_system：device 內所有 thread + host thread + 所有 peer device 的 thread。
fence「只排序，不保證可見性與原子性」；可見性需另用 volatile（繞過 L1）或 atomic 達成。

五種：Extended（cuda::atomic/atomic_ref，推薦）、Standard（cuda::std::atomic）、Built-in（__nv_atomic_*，CUDA 12.8+）、Tile、Legacy（atomic<Op>）。Legacy 僅 memory_order_relaxed、不引入 fence，scope 後綴：無=device、_block=block、_system=system。
Built-in 五級 scope：__NV_THREAD_SCOPE_THREAD/BLOCK/CLUSTER/DEVICE/SYSTEM（CLUSTER 需 sm_90+）；order 與 scope 必須是「整數字面量、不能是變數」，且僅 device、不可作用於 local memory、不可取位址。
浮點 CAS 迴圈須以 bit_cast 後的 unsigned 比較，因為 NaN != NaN：若直接用浮點比較，遇 NaN 永遠不相等會導致迴圈無法收斂而 hang。

Volta 起 thread 可獨立排程、不再保證 lock-step；mask 明確列出「應參與的 lane 集合」並在硬體執行前「強制收斂」，這是舊式無 mask 的 __shfl/__any/__ballot 被 *_sync 取代的原因（全員 0xFFFFFFFF 效率最佳）。
規則：每個呼叫 thread 的 bit 必須=1、非呼叫 thread 的 bit=0（已退出 thread 被忽略），參與者須以「相同 mask」執行；不同 thread 可用「彼此 disjoint」的 mask 並發呼叫，即使 divergent 也合法。warp *_sync 「不」提供任何 memory ordering / barrier。
__CUDA_ARCH__ 只在 device code 有定義，值 = compute_<version> × 10，取「虛擬架構」compute_80 → 800（非 sm_90）。

#pragma unroll 無引數：trip count 為常數時「完全展開」；1（或 0）：停用展開；-1（非正或 > INT_MAX）：忽略 pragma 並發警告。必須緊接在迴圈正前方。
__builtin_assume(pred) 執行期為 false 是 undefined behavior（引數有副作用則 unspecified）。
WMMA 需 CC 7.0+（bfloat16/tf32/double 需 8.0）、全 warp 協作；mptr 須 256-bit 對齊；satf=true 時 ±Inf → ±MAX_NORM、NaN → +0。
fragment 是「不透明、架構特定的 ABI 結構」，不同 link-compatible 架構（如 sm_70/sm_75）layout 不同，互傳會致錯誤/損毀；應先 store_matrix_sync 存到記憶體再以指標傳遞。

遵循 IEEE 754-2019，限制：無動態可設 rounding mode（只能用特定命名的 device intrinsic 選常數模式）、無浮點例外偵測機制、SNaN 不 signaling（當 quiet 處理）。
FMA 計算 a*b+c 只「round 一次」（非乘一次、加一次共兩次），且乘法階段保留 double-width 乘積，可防 subtractive cancellation，因此更精確。
denormal（subnormal）是 exponent 全 0、用以漸進填補最小 normal 與 0 之間空隙的值，運算較慢；-ftz=true（flush-to-zero）可沖成 0，包含於 --use_fast_math。
不論 -ftz：global 上 atomic float add「永遠 flush-to-zero」（同 add.rn.ftz.f32），shared 上「永遠支援 denormal」（同 add.rn.f32）。

intrinsic 三特性：加 __ 前綴、「只在 device code」可用、映射較少原生指令（更快但較不精確）；其行為「不受」-prec-div/-prec-sqrt/-fmad 影響，唯一例外是 -ftz=true。
非標準函式 fdividef(x,y) 是 0 ULP、等同 x/y；intrinsic __fdividef(x,y) 對 |y| ∈ [2⁻¹²⁶, 2¹²⁶] 是 2 ULP——兩者精度不同。
取整用 rint/rintf（device 上映射成「單一指令」），勿用 round/roundf（映射成「多條指令」）；trunc/ceil/floor 也都是單指令。
--use_fast_math 把一組 device Math API 函式（含 cuda::std::）自動換成對應 intrinsic（如 sinf→__sinf、x/y→__fdividef），更快但降低精度。

__pipeline_memcpy_async 只追蹤「global → shared」的非同步複製，size_and_align 只能是 4、8 或 16（且等於兩端指標的對齊）；zfill 部分補零。
tiled_partition<Size> 的 Size 須為 2 的次方且 ≤ 1024，父群大小須整除 Size；shfl_up/shfl_down/shfl_xor/ballot 僅 ≤ 32。
cg::reduce 在 CC 8.0+ 且運算為 add/min/max 或 AND/OR/XOR、且「只有 4-byte 型別」時走硬體加速，否則退回 shuffle-based。
grid.sync()（跨 grid 同步）須用 cudaLaunchCooperativeKernel（非 <<<>>>）以保證 block co-residency，僅 CC 6.0+ 支援（先查 cudaDevAttrCooperativeLaunch）。

建 stream 必須用 cudaStreamCreateWithFlags(cudaStreamNonBlocking)（cudaStreamCreate() 不可用）；event 須傳 cudaEventDisableTiming。stream handle 不可傳給 parent/child，視為建立它的 grid 私有。
device runtime「無同步函式」（cudaStreamSynchronize/Query/cudaDeviceSynchronize 皆不在子集）；要確認 stream 子 kernel 完成，改把一個 kernel 送進 cudaStreamTailLaunch。
cudaStreamFireAndForget：立即排程、不相依於先前 grid，且無 grid 能相依其完成；cudaStreamTailLaunch：排程「在自己完成之後」才啟動，達成類似 cudaDeviceSynchronize 的效果（兩者皆需 64-bit 編譯、不可定義 CUDA_FORCE_CDP1_IF_SUPPORTED）。
不支援多 GPU，只能操作當前執行的 device（但可查詢任意 device 屬性）。

因為 CUDA 的同步成本「非均勻」：執行緒距離越遠成本越高（block 內極低、跨多 GPU/CPU 很高），故在 cuda:: namespace 用 thread scope 表達這種成本，預設仍保留標準 C++ 語法語意。
由大到小：thread_scope_system ⊃ _device ⊃ _block ⊃ _thread；thread_scope_device 涵蓋「同一 device 且同一 memory synchronization domain」內的 GPU 執行緒。std::/cuda::std:: 等同以 thread_scope_system 實例化。
atomic 操作只在它「指定的 scope」內保證 atomic：scope 不是 system 時自動成立（system scope 需額外裝置屬性如 pageableMemoryAccess/concurrentManagedAccess/hostNativeAtomicSupported）。
Data race：兩個可能並行的衝突動作，至少一個「在涵蓋對方執行緒的 scope 下不是 atomic」且彼此無 happens-before → 未定義行為（UB）。

提供 parallel forward progress [intro.progress.9]，前提是 host 實作提供 concurrent FP [intro.progress.7]；parallel FP 在 thread「執行第一步前不保證開始」，一旦執行一步就升級為 concurrent FP。
傳染範圍：若屬 Cooperative Grid → 整個 grid 最終推進；否則 → 整個 thread-block cluster（蘊含整個 thread block）；其他 cluster 的 thread「不保證」推進。
CUDA 只保證「每次 API 呼叫最終 return 或確保至少一個 device thread 推進」；單次 query 不保證 producer 被排上去（parallel FP 不保證自動開始），唯有「迴圈內反覆 query」才持續觸發、保證 producer 最終推進並設旗標。
跨不同 stream 且無事件/依賴連結時，排程可任意挑推進對象（可能總挑「等待者」而餓死「設旗標者」）；須用 stream ordering 或顯式依賴串接才能可靠解開。

第五章 Technical Appendices 是「規格與語言契約」的速查層，15 篇可濃縮為下表，每列抓一個最易考的核心對照：

貫穿全章的共同精神：host 與 device 是兩個獨立執行空間，凡跨空間引用記憶體/位址/型別、或依賴 __CUDA_ARCH__ 改變介面，幾乎都是 UB；規格上限（CC、參數大小、對齊、scope）與版本/旗標契約絕不能為了便利而違反；正確性靠明確的同步/依賴/scope，不可建立在「碰巧並行」或「碰巧可見」之上。