支持的功能特性

本附录展示 cuda-oxide 的功能矩阵：每项编译器能力、运行时 API 和硬件特性， 以及其当前支持状态。数据来源于编译器/运行时源码和测试套件。

图例： Full（完整）= 已测试且正常工作，Partial（部分）= 可用但在某些方面 有已知不足（在行描述中说明），Planned（计划中）= 在路线图上， N/A = 不适用或无已识别的需求。

---

编译器：内存模型

特性	状态	描述
HMM / 统一内存管理	Full	GPU 直接读/写宿主内存，无需 cudaMemcpy。闭包中的引用捕获利用 HMM 进行宿主指针访问。需要 Turing+ GPU、Linux 6.1.24+、CUDA 12.2+。
统一结构体 ABI（无需 #[repr(C)]）	Full	设备端结构体内存布局与宿主端精确匹配。编译器查询 rustc 的实际布局，并在 LLVM IR 中用显式填充字节重现。兼容默认的 #[repr(Rust)]。
动态布局匹配	Full	编译器查询 rustc 的 fields_by_offset_order() 和字节偏移量，构建具有正确字段顺序和显式填充字节的 LLVM 结构体。独立于 LLVM 的 datalayout。

编译器：类型系统

特性	状态	描述
泛型和单态化	Full	支持带 trait 约束的泛型内核和设备函数。单态化实例从 rustc MIR 中收集。支持 const 泛型。
枚举（Option<T>、Result<T,E>、自定义枚举）	Full	完整枚举支持，包括判别式提取和 payload 访问。枚举的模式匹配正常工作。
结构体构造和字段访问	Full	结构体字面量、字段访问、按值传递和返回值。用户自定义结构体无需注解即可支持。
数组类型（[T; N]）	Full	静态数组构造、常量索引和运行时索引访问。可变数组自动提升为内存支撑。
CuSimd<T, N> SIMD 类型	Full	泛型 SIMD 寄存器类型，具名访问器（x/y/z/w）、运行时和编译期索引、to_array 转换。
ABI 标量化	Full	切片在内核边界处被标量化（&[T] -> (ptr, len)，在函数内部重新构造）。结构体和闭包按值作为单个 byval .param 传递；内部 device-to-device 调用中字段展平仍然适用。

编译器：闭包

特性	状态	描述
移动闭包（FnOnce）	Full	按值捕获的闭包。整个闭包结构体作为单个 byval 内核参数推送。move |x| x * factor 模式。
引用闭包（Fn/FnMut）	Full	非 move 闭包，按引用捕获。闭包结构体（包含宿主指针）仍作为单个 byval 参数传递；GPU 通过 HMM 读取这些指针。
宿主到设备闭包	Full	宿主端定义的闭包传递给泛型内核。已测试使用捕获系数的多项式求值。
设备内部闭包	Full	完全在设备端创建和使用的闭包，包括传递给设备函数的闭包。

编译器：控制流

特性	状态	描述
Match 表达式（整数 switch）	Full	对整数的多路匹配。生成条件分支链。
枚举 Match	Full	对 Option<T> 和自定义枚举的模式匹配。判别式提取 + payload 访问。
For 循环（range、iterator、enumerate）	Full	完整的迭代器脱糖：基于 range、slice.iter()、enumerate()、嵌套循环、break、continue。
While 循环 / If-Else	Full	基本控制流完全支持。
Break 和 Continue	Full	for/while 循环中的 break 和 continue，包括提前退出。

编译器：算术和类型转换

特性	状态	描述
64 位算术运算	Full	完整的 64 位整数算术运算，包括移位、按位操作和描述符字段打包。
类型转换（所有类型）	Full	IntToInt、IntToFloat、FloatToInt、FloatToFloat、Transmute（bitcast）、PtrToPtr、PtrToInt、IntToPtr、指针强制转换。

编译器：互操作

特性	状态	描述
双向 LTOIR 支持	Full	Rust 内核调用 CUDA C++ 设备函数，同时 C++ 调用 Rust 设备函数。通过 NVVM IR → libNVVM → LTOIR → nvJitLink。
设备端 FFI（extern "C"）	Full	#[device] extern "C" { fn ... } 声明用于外部 LTOIR 函数。已验证 CUB/CCCL 集成。
MathDx FFI（cuFFTDx / cuBLASDx）	Full	cuFFTDx（8/16/32 点线程级 FFT）、cuBLASDx（32x32x32 块级 GEMM）通过 LTOIR。
跨 Crate 内核	Full	在库 crate 中定义的内核和设备函数，在二进制 crate 使用点进行单态化。

编译器：函数

特性	状态	描述
#[kernel] 属性	Full	将函数标记为 GPU 内核入口点（ptx_kernel 调用约定）。每个文件可包含多个内核。
#[device] 辅助函数	Full	可从内核调用的设备端辅助函数。由 llc 激进内联。
独立的 #[device] 函数	Full	在没有任何内核的情况下编译的设备函数。为 C++ 消费生成干净的导出名。
多内核模块	Full	单个源文件中的多个 #[kernel] 函数编译为单个 PTX 模块。

编译器：编译管线

特性	状态	描述
统一单源编译	Full	宿主和设备代码在同一文件中。自定义 rustc 代码生成后端拦截代码生成。无需 #[cfg]。
PTX 输出	Full	默认输出：Rust MIR → dialect-mir → mem2reg → dialect-llvm → LLVM IR → llc → PTX。目标架构 sm_80 至 sm_100a。
NVVM IR 输出	Full	为 libNVVM 消费提供的替代输出，包含 NVVM 元数据。
LTOIR 链接	Full	通过 libNVVM 和 nvJitLink 进行设备端 LTO。
浮点数学内建函数（libdevice）	Full	Rust f32/f64 数学方法（sin、cos、exp、pow、sqrt 等）降层为 CUDA libdevice（__nv_*）。cuda-oxide 自动检测 libdevice 使用并生成 NVVM IR；cuda_host::load_kernel_module（同步）和 cuda_host::load_kernel_module_async（异步）在运行时通过 libNVVM + nvJitLink 构建 cubin。
管线检查	Full	cargo oxide pipeline <example> 显示每个编译阶段的 IR。
cuda-gdb 调试支持	Full	带调试信息构建并使用 cuda-gdb 启动。breakpoint() 内建函数用于程序化断点。

---

运行时库：安全性

特性	状态	描述
DisjointSlice<T, IndexSpace>	Full	带边界检查的并行写入输出切片。get_mut 和 get_mut_indexed 返回 Option<&mut T>。IndexSpace 类型参数在编译期拒绝不匹配的 2D 步长。
ThreadIndex<'kernel, IndexSpace>	Full	不透明 witness，仅可由受信任的索引函数构造。!Send + !Sync + !Copy + !Clone，'kernel 作用域限定——不可跨线程传递，不可存活超过内核体。
ManagedBarrier 类型状态	Full	编译期屏障生命周期：Uninit → Ready → Invalidated。无效的状态转换是编译错误。

运行时库：原子操作

特性	状态	描述
Device 作用域原子操作	Full	DeviceAtomic{U32,I32,U64,I64,F32,F64} 使用 .gpu 作用域。支持全部 5 种排序。
Block 作用域原子操作	Full	BlockAtomic{U32,I32,U64,I64,F32,F64} 使用 .cta 作用域。
System 作用域原子操作	Full	SystemAtomic{U32,I32,U64,I64,F32,F64} 使用 .sys 作用域。用于 CPU-GPU 共享数据。
core::sync::atomic 支持	Full	标准库原子类型降层为 PTX atom.sys 指令。

运行时库：共享内存

特性	状态	描述
静态共享内存	Full	SharedArray<T, N, ALIGN>——编译期固定大小，块级作用域。可选对齐，最高 256B。
动态共享内存	Full	DynamicSharedArray<T, ALIGN>——运行时大小，通过 LaunchConfig::shared_mem_bytes 设置。
分布式共享内存（DSMEM）	Full	在集群内直接访问其他块的共享内存。map_shared_rank() 用于地址映射。sm_90+。

运行时库：线程和同步

特性	状态	描述
线程/块/网格内建函数	Full	threadIdx、blockIdx、blockDim、gridDim。index_1d() 和 index_2d::<S>()（常量步长）是类型安全的；index_2d_runtime(s) 是步长仅在 launch 时可知情况下的 unsafe 紧急出口。参见 安全模型。
块级同步	Full	sync_threads()——线程块屏障。
异步屏障（mbarrier）	Full	Hopper+ 的硬件异步屏障：init、arrive、test_wait、try_wait、inval。
集群同步	Full	针对集群中所有块的 cluster_sync()。sm_90+。
Fence 操作	Full	用于 TMA 可见性的 fence_proxy_async_shared_cta()、nanosleep(ns)。

运行时库：Warp

特性	状态	描述
Warp Shuffle 操作	Full	shuffle、shuffle_xor、shuffle_down、shuffle_up，支持 i32 和 f32。
Warp Vote 操作	Full	all(pred)、any(pred)、ballot(pred) → 位掩码。
Lane/Warp ID	Full	lane_id()（0–31）、warp_id()。直接寄存器读取。

运行时库：协作组

特性	状态	描述
类型化 Group 句柄	Full	Grid、Cluster、ThreadBlock、WarpTile<N>（N ∈ {1,2,4,8,16,32}）、CoalescedThreads。
Group 通用 API	Full	每个 group 句柄上的 size()、thread_rank()、sync()。
Warp Tile 分区	Full	ThreadBlock::tiled_partition::<N>() 划分出子 warp WarpTile<N>。coalesced_threads() 具象化活跃 lane 组。
Warp 集合操作	Full	ballot、all、any、shfl、shfl_xor、shfl_down、shfl_up（i32 和 f32）；match_any / match_all（i32 和 i64）；active_mask。
Warp 归约 / 扫描	Full	warp_reduce、warp_scan（包含式）。对 u32/i32/f32 的 Sum/Min/Max；对 u32 的 BitAnd/BitOr/BitXor。
Block 归约 / 扫描	Full	block_reduce、block_scan（包含式）。通过 const 泛型 NUM_WARPS 参数化；与 warp 变体相同的操作/类型矩阵；使用 __shared__ 暂存区。
协作内核 Launch	Full	cuda_launch! { cooperative: true, ... } 启用 Grid::sync() 进行网格级屏障。

运行时库：调试

特性	状态	描述
gpu_printf! 宏	Full	格式化 GPU 输出，支持完整的格式说明符。降层为 vprintf。
gpu_assert! 宏	Full	GPU 运行时断言。条件为假时调用 trap()。
调试内建函数	Full	clock()、clock64()、trap()、breakpoint()、prof_trigger::<N>()。

运行时库：内核 Launch

特性	状态	描述
#[cuda_module] 类型化 Launch	Full	嵌入式模块加载，带类型化同步/异步 launch 方法。
cuda_launch! 宏	Full	较低级别的 launch，显式模块加载和封装器。
#[launch_bounds]	Full	Occupancy 提示：每块最大线程数、每 SM 最小块数。
#[cluster_launch]	Full	编译期集群维度。在 PTX 中生成 .reqnctapercluster。

运行时库：TMA

特性	状态	描述
TMA 批量张量拷贝（1D–5D）	Full	cp_async_bulk_tensor_{1..5}d_g2s。128 字节 TMA 描述符。sm_90+。
TMA 多播	Full	单次 TMA 加载广播到集群中的所有 CTA。完全多播需 sm_100a。
TMA Commit/Wait 组	Full	cp_async_bulk_commit_group、cp_async_bulk_wait_group 用于异步完成跟踪。

---

尚未实现

特性	状态	备注
内联汇编（asm! 宏）	Planned	变通方案：使用内置内建函数或向 cuda-device 添加新的内建函数。
FP8 / MX 数据类型	Planned	Blackwell 路线图项目。无架构限制。
动态分发（dyn Trait）	N/A	使用泛型和静态分发。尚未发现对此的实际需求。
堆分配（Box、Vec）	N/A	CUDA 有设备端堆（内核中的 malloc/free），编译器允许 alloc crate 通过——但目前未连接设备端 #[global_allocator]。即便连接了，设备端 malloc 也极其缓慢（串行化、碎片化、非合并）。请使用切片和 SharedArray。
String / format_args!	N/A	使用 gpu_printf! 进行格式化输出。
Panic / Unwinding	N/A	Panic 路径存在于 MIR 中，但编译器会剥离 core::panicking::* 和所有 unwind 边。GPU 硬件可以支持 unwinding（绝对分支 + Volta 后的每线程调用栈跟踪），但 CUDA 工具链（nvcc/ptxas）目前不提供——没有 landing pad 能存活到 PTX。如果运行时到达 panic 路径，GPU 会 trap（等同于 panic=abort）。NVIDIA 有一个正在进行的项目，旨在将 C++ 异常支持添加到 CUDA 中以用于汽车安全领域；当前的 cuda-oxide 设计与该工作前向兼容。目前使用 gpu_assert!() + trap() 进行显式运行时检查。
标准库（std/alloc）	N/A	std 被禁止。alloc 被收集器允许，但没有后备分配器。只有 core 完全功能可用。Option、Result、迭代器均可正常工作。
纹理内存	N/A	考虑到 Hopper+ 上 TMA 的可用性，优先级较低。