术语表#
以下术语按 cuda-oxide 项目和本书中的用法定义。
ABI 标量化（ABI Scalarization）#
将复合类型分解为宿主端和设备端一致认可的形状的过程。切片在所有调用边界处展平：
&[T] 变为 (ptr, len) 对，在函数内部通过 insertvalue/extractvalue 重新构造。
结构体和按值传递的闭包在内部 device-to-device 调用中以同样方式展平，但在
内核边界处，它们作为单个 byval .param 传递（一个宿主数据包槽位）
——这与宿主 launcher 推送的内容匹配，避免了展平的设备声明与单槽宿主数据包之间的不匹配。
块（Block / Thread Block）#
一组在同一 Streaming Multiprocessor 上执行并共享对共享内存访问的线程。
块内线程可通过 sync_threads() 同步。由 blockIdx 标识，由 blockDim 指定大小。
集群（Cluster / Thread Block Cluster）#
Hopper+（sm_90）特性：一组最多 16 个线程块，保证在同一 GPC 上协同调度。
集群中的块可通过分布式共享内存（DSMEM）访问彼此的共享内存，
并通过 cluster_sync() 同步。通过 #[cluster_launch] 配置。
代码生成后端（Codegen Backend）#
rustc_codegen_cuda crate——作为 dylib 加载的自定义 rustc 后端。它在编译期间
拦截 MIR，通过 dialect-mir → mem2reg → dialect-llvm → LLVM IR → PTX
进行降层，并将 PTX 与常规宿主二进制文件一起输出。
cuda-async#
异步执行层。提供 DeviceOperation（惰性 GPU 工作描述）、DeviceFuture
（绑定到流的执行）和 DeviceBox<T>（设备端拥有的内存）。通过
zip!、and_then 和 value() 组合工作。
cuda-device#
#![no_std] 设备端 crate，提供所有 GPU 内建函数和类型：线程标识、共享内存、
warp 原语、屏障、TMA、张量核心、原子操作和调试设施。
cuda-core#
围绕 CUDA Driver API 的安全 RAII 封装器：CudaContext、CudaStream、
DeviceBuffer<T> 和模块加载。在宿主端处理 GPU 上下文和内存管理。
DeviceOperation#
惰性、可组合的 GPU 工作描述（分配、内核 launch 或数据传输）。
在调用 .sync() 或 .await 之前不会执行。可通过
zip!（并行）和 and_then（顺序）组合。
DisjointSlice<T, IndexSpace>#
用于内核的安全可变输出抽象。仅接受 IndexSpace 与其自身类型参数匹配的
ThreadIndex，提供带边界检查的 Option<&mut T> 返回值。通过构造防止数据竞争
——每个线程只能写入自己的元素。get_mut_indexed() 快捷方法在一次调用中
生成 witness 并将其解析为可变引用。
网格（Grid）#
内核 launch 中线程的顶层组织。网格是一个 3 维线程块数组，由 gridDim 指定大小。
总线程数为 gridDim × blockDim。
HMM（异构内存管理 / Heterogeneous Memory Management）#
Linux 内核（6.1.24+）特性，允许 GPU 通过页错误直接访问宿主内存，无需显式的
cudaMemcpy。cuda-oxide 利用 HMM 实现闭包中的引用捕获
——GPU 透明地读取宿主地址。也称为统一内存管理（UMM）。
Lane#
warp 内的单个线程。Lane 编号为 0–31，由 warp::lane_id() 标识。
Warp shuffle 和 vote 操作在 lane 之间通信，无需共享内存或屏障。
LTOIR（链接时优化 IR / Link-Time Optimized IR）#
用于设备端链接时优化的中间表示。cuda-oxide 为 libNVVM 生成 NVVM IR，
libNVVM 可将其编译为 LTOIR，以便使用 nvJitLink 将 Rust 设备代码与
CUDA C++ 设备代码链接。
ManagedBarrier<State, Kind, ID>#
Hopper+ 上用于异步操作的类型状态屏障。在编译期跟踪其生命周期
（Uninit → Ready → Invalidated）和用途（TmaBarrier、
MmaBarrier、GeneralBarrier）。无效的状态转换是编译错误。
单态化（Monomorphization）#
Rust 编译器为每个使用的具体类型生成泛型函数特化副本的过程。cuda-oxide
完全支持设备端的单态化——scale::<f32> 和 scale::<f64> 各生成独立的 PTX 函数。
Pliron#
用 Rust 编写的 IR 框架，受 MLIR 启发，用作 cuda-oxide 编译管线中的中间表示。 MIR 被导入到 Pliron IR 中，通过方言 pass 进行转换，并导出为 LLVM IR。
PTX（并行线程执行 / Parallel Thread Execution）#
NVIDIA 的 GPU 内核底层虚拟 ISA。cuda-oxide 编译器生成 PTX 作为主要输出， CUDA 驱动在加载时将其 JIT 编译为原生 GPU 机器码（SASS）。
SM（流式多处理器 / Streaming Multiprocessor）#
NVIDIA GPU 上的主要处理单元。每个 SM 拥有自己的寄存器、共享内存、warp 调度器和执行管线（包括张量核心）。线程块由硬件调度器调度到 SM 上。
sync_threads()#
块级屏障：线程块中的所有线程必须到达此处后，任何线程才能继续执行。
等价于 CUDA C++ 中的 __syncthreads()。降层为 llvm.nvvm.barrier0()。
张量核心（Tensor Cores）#
专用的矩阵乘加硬件单元。WGMMA（Hopper, sm_90）以 warpgroup 粒度从共享内存 操作。tcgen05（Blackwell, sm_100+）使用单线程发起和专用张量内存（TMEM）。
ThreadIndex<'kernel, IndexSpace>#
不透明 witness，只能由受信任的索引函数构造。三种形式：
thread::index_1d() -> ThreadIndex<'_, Index1D>。始终返回 witness； 每个线程无条件唯一（threadIdx.x < blockDim.x由硬件保证）。thread::index_2d::<S>() -> Option<ThreadIndex<'_, Index2D<S>>>。 行步长是 const 泛型，因此DisjointSlice<T, Index2D<S>>只接受具有匹配S的 witness——混合步长是类型错误。unsafe thread::index_2d_runtime(s) -> Option<ThreadIndex<'_, Runtime2DIndex>>。 当步长仅在 launch 时可知时的紧急出口。unsafe是契约：每个将Runtime2DIndex送入同一个DisjointSlice的线程必须使用相同的s。
Witness 是 !Send + !Sync + !Copy + !Clone 且 'kernel 作用域限定，
因此线程无法通过共享内存传递它，也不能在内核体之后继续存活。
TMA（张量内存加速器 / Tensor Memory Accelerator）#
Hopper+ 硬件单元，用于全局内存和共享内存之间的异步批量拷贝。通过
TmaDescriptor（在宿主端构建的 128 字节不透明描述符）和
cp_async_bulk_tensor_* 内建函数操作。完成情况由 ManagedBarrier 跟踪。
TmemGuard<State, N_COLS>#
Blackwell 张量内存（TMEM）的类型状态封装器——用于 tcgen05 MMA 操作的专用累加器
存储。管理 TMEM 生命周期：TmemUninit → TmemReady → TmemDeallocated。
无效转换是编译错误。
Warp#
一组 32 个线程，在 SM 上以锁步方式执行指令。Warp 是最小调度单元。 Warp 级别操作（shuffle、vote）在约 1 个周期内在 lane 之间交换数据， 无需共享内存或同步屏障。
Warpgroup#
四个连续的 warp（128 线程），在 Hopper 上为 WGMMA 操作协同运行。 Warpgroup 是 warpgroup 级别矩阵乘加的发起单元。
WGMMA（Warpgroup 矩阵乘加 / Warpgroup Matrix Multiply-Accumulate）#
Hopper（sm_90）的张量核心指令集。四个 warp 协同从共享内存操作数向寄存器
累加器发起 MMA。通过 ManagedBarrier 进行异步执行和 commit/wait。