fn 已弃用;改用 def。 Mojo 现在对 fn 的使用发出编译器警告;这将在下一个版本中变为编译错误。这完成了自 v0.26.2 开始的 def/fn 统一:def 现在是 Mojo 的标准函数声明关键字,具有 fn 原本的不可抛出语义。请参阅语言变更。
统一的闭包。 此版本延续了早期版本开始的闭包统一工作:无状态闭包自动提升为顶层函数(可作为 FFI 回调传递),支持 ref 捕获约定,默认捕获约定可与显式捕获列表组合,且新增的 thin 函数效果声明了不携带捕获状态的纯函数指针类型。请参阅语言增强。
UnsafePointer 默认非空。 默认的 null 构造函数和 __bool__() 方法已弃用,UnsafePointer 不再遵循 Defaultable 或 Boolable。使用 Optional[UnsafePointer[...]] 表达可空性,它与 UnsafePointer 共享内存布局(空地址是 None 的 niche),因此可空指针保持零开销和 FFI 安全。请参阅指针和内存。
集合默认开启边界检查。 所有标准库集合已移除负索引——x[-1] 现在是编译时错误;改用 x[len(x) - 1]——且 CPU 上的所有集合默认开启边界检查。越界访问会报告用户的调用位置。GPU 上为性能默认仍关闭边界检查;使用 mojo build -D ASSERT=all 启用。请参阅集合与迭代器。
NDBuffer 已移除。 NDBuffer 已从标准库中完全移除。迁移至 TileTensor。请参阅集合与迭代器。
扩展的 GPU 硬件支持。 Apple Metal 成为功能更强的 Mojo 目标——print() 可用,支持动态线程组内存(external_memory[]()),Apple M5 MMA 内建函数支持硬件矩阵乘累加,且 Apple GPU 目标默认优先使用 metal4 特性。新增对 AMD MI250X 和 NVIDIA B300 (sm_103a) 加速器的支持。请参阅 GPU 编程。
GPU 原始 ID 访问器迁移 UInt → Int。 thread_idx、block_idx、block_dim、grid_dim、global_idx、lane_id、warp_id 以及集群访问器现在返回 Int,这是标准化使用 Int 表示大小和偏移的更广泛迁移的一部分。临时的 *_uint 别名提供了迁移路径;它们最终将被弃用并移除。请参阅 GPU 编程。
CPU DeviceContext 扩展。 DeviceContext(api="cpu") 现在是一个用于 CPU 工作的流排序执行上下文,为 NUMA 感知 CPU 调度铺平了道路。新增 enqueue_cpu_function() 和 enqueue_cpu_range() 可将主机函数和并行范围以流排序的方式相对于周边工作入队。请参阅 GPU 编程。
String 和 StringSlice 中的字形簇支持。 新增基于 UAX #29 的字形簇分割,提供 graphemes()、count_graphemes()、[grapheme=...] 切片语法以及反向迭代。正确处理组合标记、emoji ZWJ 序列、旗帜 emoji、韩语音节和其他多码点簇。请参阅字符串与文本。
类型精化。 编译器现在根据 where 子句、comptime if 语句和 comptime assert 语句缩窄类型,由 conforms_to() 表达式驱动。这使得 trait_downcast 在常见情况下变得不再必要——Mojo 能识别出类型在精化作用域内满足某个 trait,并允许你直接调用其 trait 方法。请参阅语言增强。
统一的反射 API。 std.reflection 中新增的 reflect[T]() 入口点返回一个 Reflected[T] 句柄,取代了 struct_field_* 系列自由函数以及 get_type_name / get_base_type_name。reflect 通过预导入自动可用。旧版自由函数和 ReflectedType[T] 包装器现已 @deprecated。请参阅其他库变更。
新增基于编译时假设的类型精化,使 Mojo 能够从 where 子句、comptime if 语句和 comptime assert 语句缩窄类型。作用域内的精化由 conforms_to() 表达式驱动。
之前:
def __contains__(self, value: Self.T) -> Bool where conforms_to(Self.T, Equatable):
for item in self:
if trait_downcast[Equatable](item.md) == trait_downcast[Equatable](value.md):
return True
return False
之后:
def __contains__(self, value: Self.T) -> Bool where conforms_to(Self.T, Equatable):
for item in self:
if item == value:
return True
return False
统一的闭包改进。此版本延续了早期版本开始的闭包统一工作:无状态闭包自动提升,支持 ref 捕获约定,默认捕获约定可与显式捕获列表组合,且新增的 thin 函数效果声明了不携带捕获状态的纯函数指针类型。
def main() raises:
var a, b, c, d = 1, 2, 3, 4
var x = "hello"
# 旧式闭包:无捕获列表。不能捕获变量。
def hello():
print("hi")
# 无捕获的统一闭包(无状态)。无状态闭包提升为顶层函数,可作为 FFI 回调传递。
def add_one(n: Int) {} -> Int:
return n + 1
# 带有显式捕获和默认捕获约定的统一闭包:
def my_fn() {mut a, b, c^, read}:
# 捕获:
# `a` 以可变引用捕获
# `b` 以不可变引用捕获
# `c` 以移动捕获
# `d` 以不可变引用捕获(默认的 `read` 约定)
use(a, b, c, d)
# 以 ref 方式捕获 `x` 的统一闭包(携带原始可变性参数):
def show_x() {ref x}:
print(x)
# 函数效果置于捕获列表之前。调用上下文必须处理从 `raises` 闭包中抛出的错误。
def fallible() raises {}:
raise Error("nope")
# 闭包与普通函数一样调用:
hello()
print(add_one(41))
my_fn()
show_x()
try:
fallible()
except e:
print(e)
# `thin` 函数效果声明一个不携带捕获状态的纯函数指针类型。无状态闭包和顶层函数与 `thin` 函数指针兼容:
var fn_ptr: def(Int) thin -> Int = add_one
print(fn_ptr(99))
新增 abi("C") 作为函数效果,用于在函数定义和函数指针类型上声明 C 调用约定。标记了 abi("C") 的函数对结构体参数和返回值使用平台 C ABI(System V x86-64 / ARM64 AAPCS),从而能够与 C 库安全互操作。DLHandle.get_function() 现在强制要求类型参数携带 abi("C"),防止加载 C 符号时出现无声的 ABI 不匹配。
# C-ABI 函数定义(可安全用作 C 代码的回调)
def add(a: Int32, b: Int32) abi("C") -> Int32:
return a + b
# C-ABI 函数指针类型(可安全用于 DLHandle.get_function)
var f = handle.get_function[def(Float64) abi("C") -> Float64]("sqrt")
新增对条件 RegisterPassable 遵循的支持。
三元 if/else 表达式现在在显而易见时将每个元素强制转换为上下文类型。例如,以下代码现在能正常工作,而不再产生不兼容元类型错误:
comptime some_type: Movable = Int if cond else String
当被调用方接受兼容的可变参数列表或包时,可变参数列表和包可通过 *pack 通过运行时调用转发。
def callee[*Ts: Writable](*args: *Ts):
comptime for i in range(args.__len__()):
print(args[i])
def forwarder[*Ts: Writable](*args: *Ts):
callee(*args)
forwarder(1, "hello", 3.14) # 每行打印一个值
异构可变参数包现在可使用 SomeTypeList 助手指定。以下两种写法等价:
def foo[*arg_types: Copyable](*args: *arg_types) -> Int: ...
def foo(*args: *SomeTypeList[Copyable]) -> Int: ...
字符串字面量现在支持 \uXXXX 和 \UXXXXXXXX Unicode 转义序列,与 Python 一致。生成的码点以 UTF-8 存储。无效码点和代理对在解析时拒绝。
T-字符串现在可用于 comptime assert 消息中:
def foo[i: Int]():
comptime assert i > 5, t"expected i > 5, got {i}"
用于函数声明的 fn 关键字已弃用。Mojo 现在对 fn 的使用发出编译器警告;这将在下一个版本中变为编译错误。改用 def。
unified 关键字已移除;通过在函数签名后使用显式捕获列表 {...} 来指定统一闭包语义。空捕获列表 {} 表示无捕获的统一闭包;不带任何捕获列表的闭包属于旧式。Mojo 现在还警告在函数指针类型中省略 thin 效果;显式指定 thin 可消除警告。
from pkg import ... 形式的导入语句不再让 pkg 对模块可用。
移除了对不同长度或类型元组进行比较的支持。此类比较(例如 (1, 2) != (4, 5, 6))现在被类型系统静态拒绝,而不再默默返回不相等。
可变参数列表现在是 ParameterList 和 TypeList 而不是 !kgen.param_list,因此可以像普通类型一样使用:
def callee[*values: Int]():
var v = 0
for i in range(len(values)):
v += values[i]
for elt in values:
v += elt
每个 Mojo 函数现在拥有自己唯一的函数字面量类型。即使签名完全相同,两个单独定义的函数也不能通过其字面量类型互换使用;使用函数指针类型(例如 def(Int) thin -> Int)来对它们进行抽象。
A if comptime(C) else B 现在跳过对死分支的展开,将三元表达式视为类似于 comptime if C: A else: B 的编译时求值契约。
当与无条件的 ImplicitlyDestructible 遵循配对时,@explicit_destroy 现在在解析时报错;它仅在条件(受 where 子句约束的)遵循上保持有效。
NDBuffer 已完全移除。迁移至 TileTensor。
所有标准库集合(List、Span、InlineArray、String、StringSlice、LinkedList、Deque、IntTuple)已移除负索引,以便默认启用低开销的 CPU 边界检查。使用负 IntLiteral 进行索引现在会触发编译时错误:
constraint failed: negative indexing is not supported, use e.g. `x[len(x) - 1]` instead
将任何 x[-1] 更新为 x[len(x) - 1]。
CPU 上的所有集合现在默认启用边界检查。越界访问会报告用户的调用位置:
def main():
var x = [1, 2, 3]
print(x[3])
At: /tmp/main.mojo:3:12: Assert Error: index 3 is out of bounds, valid range is 0 to 2
GPU 上为性能默认仍关闭边界检查。使用 mojo build -D ASSERT=all 在 GPU 上启用边界检查;使用 -D ASSERT=none 禁用所有断言,包括 CPU 边界检查。
接受不同类型参数或接受 Intable/IntableRaising 但不是 Indexer 参数的 range() 重载已移除。调用方应传递一致的整数参数类型。
在迭代模块中新增 IterableOwned trait。遵循 IterableOwned 的类型实现 __iter__(var self),该函数消耗集合并返回拥有底层元素的迭代器。List、Optional、Deque、LinkedList、Dict、Set、Counter 和 InlineArray 现在遵循此 trait;Span 在 T: Copyable 时有条件地遵循,其自有迭代器按值生成副本。
迭代器适配器(enumerate()、zip()、map()、peekable()、take_while()、drop_while()、product()、cycle()、count()、repeat())现在遵循 IterableOwned。为 enumerate()、zip()、map()、peekable()、take_while()、drop_while()、product() 和 cycle() 新增了消耗输入可迭代对象的自有重载。
为 Optional 新增 map() 和 and_then() 方法。map() 将函数应用于包含的值(返回 Optional[To]);and_then() 对自身返回 Optional 的操作进行 flat-map。
var o = Optional[Int](42)
def closure(n: Int) {} -> String:
return String(n + 1)
var mapped: Optional[String] = o.map[To=String](closure)
print(mapped) # Optional("43")
新增 Optional.destroy_with(destroy_func),它使用调用者提供的析构函数就地销毁 Optional[T]。这使得 Optional 能够持有不是 ImplicitlyDestructible 的元素类型(例如标记了 @explicit_destroy 的类型),镜像了 Variant.destroy_with()。两个 destroy_with() 方法现在除了接受纯函数引用外,还接受捕获局部状态的闭包。Variant.destroy_with() 调用方现在必须显式传递销毁的类型(例如 v^.destroy_with[Int](destroy_func)),因为 T 不再能从闭包类型推断。
为 Span 新增了泛型 __contains__() 方法,适用于遵循 Equatable 的任何元素类型,而不仅仅是 Scalar 类型。
assert_raises() 现在捕获自定义 Writable 错误类型,而不仅仅是 Error。
新增对 AMD MI250X 加速器的支持。
扩展的 Apple 芯片 GPU 支持。Apple Metal GPU 现在是一个功能更强的 Mojo 目标。
print() 和 _printf() 现在在 Apple Metal GPU 上可用。输出通过 Metal os_log 路径分块,使用匹配 Metal 硬件约束的 Float32 专用格式化器。_printf() 目前仅发出格式字符串(不插值参数);|x| < 1e-7 被截断为 0.0。
external_memory[]()(动态线程组内存)现在在 Apple 芯片上受支持,因此现有的使用 external_memory[]() 的 GPU 内核可保持不变地工作。
std.gpu.compute.arch.mma_apple 中的 Apple M5 MMA 内建函数(apple_mma_load()、apple_mma_store()、_mma_apple())支持在 Apple GPU 上进行硬件矩阵乘累加。
在 std.sys 中新增 CompilationTarget.is_apple_m5() 用于在编译时检测 Apple M5 目标;is_apple_silicon() 现在在其检查中包含 M5。
当工具链支持时,Apple GPU 目标现在默认优先使用 metal4 特性,自动将 -metal4 附加到架构字符串,而不再需要显式选择 m5-metal4。
原子排序:release 排序在 Metal 上不受支持。Apple GPU 目标现在默认使用 monotonic(松弛)原子排序。
浮点宽度:编译器现在拒绝 Apple GPU 目标上宽度超过 32 位的浮点类型(Float64/Float80/Float128),因为 Metal 仅支持 Float16 和 Float32。
GPU 设备 API:
新增对 NVIDIA B300 (sm_103a) 加速器的支持。std.sys.info 和 std.gpu.host.info 中的新助手可识别 B300 目标,以便在 Blackwell B300 上正确调度内核。
新增 DeviceStream.enqueue_host_func(func, user_data),为 Mojo 内核和自定义算子暴露 cuLaunchHostFunc 原语。接受 thin def(OpaquePointer[MutAnyOrigin]) -> None 回调和 opaque user_data 指针。目前仅 CUDA 可用;非 CUDA 后端抛出错误。
DeviceContext 初始化现在运行自动 GPU 健康检查,检测硬件降频、不可纠正的 ECC 错误和僵尸 VRAM,并在 GPU 不健康时以可操作的错误消息使设备创建失败。新增 DeviceContext.run_healthcheck() 以显式重新调用检查。设置 MODULAR_DEVICE_CONTEXT_DISABLE_HEALTHCHECK=true 禁用。
优化了 GPU elementwise() 索引计算和调度,采用 use_32bit 快速路径、4 倍展开的网格步进处理、warp 对齐的块大小和 SM100+ 单 tile 路由。
GPU 原始 ID 访问器(thread_idx、block_idx、block_dim、grid_dim、global_idx、lane_id、warp_id、cluster_dim、cluster_idx 和 block_id_in_cluster)已从 UInt 迁移到 Int。
这是标准化使用 Int 类型表示 Mojo 中所有大小和偏移的更广泛迁移的一部分。作为同一次迁移中的相关步骤,TensorCore.load_a() 和 TensorCore.load_b() 现在也接受 Int 参数而不是 UInt。
为提供渐进的迁移路径,七个非集群访问器临时提供 *_uint 别名:
| 访问器 | 旧版 UInt 别名 |
|---|---|
| thread_idx | thread_idx_uint |
| block_idx | block_idx_uint |
| block_dim | block_dim_uint |
| grid_dim | grid_dim_uint |
| global_idx | global_idx_uint |
| lane_id | lane_id_uint |
| warp_id | warp_id_uint |
三个集群访问器(cluster_dim、cluster_idx、block_id_in_cluster)直接迁移,没有 *_uint 别名,因为它们的使用有限。
代码可以通过对 *_uint 别名使用重命名导入来保持其先前的 UInt 行为:
- from std.gpu import thread_idx
+ from std.gpu import thread_idx_uint as thread_idx
在分阶段迁移期间短暂存在的临时 *_int 访问器(用于帮助前向兼容)已移除;使用无前缀的访问器(它们现在默认返回 Int)。*_uint 别名最终将被弃用并移除。
CPU DeviceContext 扩展。DeviceContext(api="cpu") 现在可用作 CPU 工作的流排序执行上下文,为 NUMA 感知 CPU 调度铺平道路。
新增 DeviceContext.enqueue_cpu_function() 和 DeviceContext.enqueue_cpu_range(),用于在 CPU DeviceContext 实例上对流排序执行主机函数。enqueue_cpu_function() 入队单个主机函数;enqueue_cpu_range() 入队一个并行范围,其任务并发运行,但相对于周边工作是流排序的。尚不支持参数传递。
parallelize()、parallelize_over_rows()(在 std.algorithm.backend.cpu.parallelize 中)以及 std.algorithm.functional 中的 elementwise() 重载现在接受可选的尾随参数 ctx: Optional[DeviceContext] = None。当提供时,上下文被转发给 sync_parallelize();否则行为不变。
新增接受 CPU DeviceContext 并返回该特定上下文线程池大小的 parallelism_level() 重载,支持 NUMA 特定的内省。
AMD GPU 内建函数:
在 std.gpu.intrinsics 中新增 ds_read_tr8_b64() AMD GPU 内建函数,通过 llvm.amdgcn.ds.read.tr8.b64 执行 8 位元素的 64 位 LDS 转置加载。在 AMD CDNA4+ GPU 上受支持。
为 readfirstlane() 新增 Scalar[dtype] 重载,使调用方不再需要绕弯通过位转换来在 AMD GPU wavefront 上广播非 Int32 标量值。
std.gpu.intrinsics 中的 AMDBufferResource.load_to_lds() 现在 lower 为 AMDGPU 缓冲区加载到 LDS 内建函数的 .ptr. 形式,修复了一个跨步布局回归。新增的 async_copies: Bool = False 参数选择附加 amdgpu.AsyncCopies 别名作用域到加载操作,支持 LLVM vmcnt 松弛。
为 GPU warp_id()(及相关 ID 访问器)新增 broadcast=True 参数,使调用方可避免手动的 warp.broadcast(warp_id()) 模式。
用于 TileTensor 数据移动的 tile_io 模块。新增 tile_io 模块,提供用于在内存层次(DRAM/SRAM)之间移动数据的 TileTensor 复制 trait 和复制实用工具。该模块包含:
GenericToSharedAsyncTileCopier,它通过 NVIDIA 的 cp.async 将 TileTensor 从通用内存移动到共享内存。在 AMD 和 Apple GPU 上,底层的 async_copy() 回退到同步加载/存储。
GenericToSharedAsyncTileCopier 上可选的 swizzle: Swizzle 参数,镜像 LocalToSharedTileCopier 中的交错写入路径。
GenericToSharedAsyncTileCopier 上的 masked: Bool = False 参数。启用时,越界向量会收到零字节拷贝并零填充,匹配 LayoutTensor.copy_from_async[is_masked=True, fill=Fill.ZERO]。
用于抽象复制器遵循的 AsyncTileCopier trait。
用于稀疏 2D 张量加载的 TMA gather4。在 SM100 (Blackwell) 上新增 TMA gather4 操作,用于在单条 TMA 指令中加载来自 2D 张量的 4 个非连续行,表面化为 std.gpu.memory 中的 cp_async_bulk_tensor_2d_gather4() 内建函数,并与 TMATensorTile 集成。该 API 支持:
适用于 SM90+ NVIDIA GPU 的 1D TMA 指令。在 std.gpu.memory 中新增 1D TMA(张量内存加速器)指令支持。1D TMA 拷贝不需要在主机上预分配张量映射对象,提供了比现有 2D–5D TMA 路径更大的灵活性。新函数:cp_async_bulk_shared_cluster_global()、cp_async_bulk_global_shared_cta()、cp_async_bulk_prefetch() 和 cp_async_bulk_reduce_global_shared_cta()。
在剖析器中可读的 GPU 内核名称。标准库和整个 MAX 内核(elementwise、GEMV、多级 matmul、attention、卷积、MoE、归一化、量化、BMM、分组 matmul、SM100 matmul、AMD matmul、通信和采样)中的 GPU 内核现在在 Nsight Systems 等剖析器跟踪中暴露人类可读的名称,取代先前重整的 KGEN 符号。
在 DeviceContext、DeviceBuffer 和 HostBuffer 上为 enqueue_copy()、enqueue_copy_from() 和 enqueue_copy_to() 新增基于 Span 的重载,为主机到设备的内存传输提供比原始 UnsafePointer 更安全的替代方案。
从整个标准库中移除了 trait_downcast() 和 trait_downcast_var(),取而代之的是类型精化(参见语言增强)。公共 API 不变。
external_call() 的 return_type 要求已从 TrivialRegisterPassable 放宽到 RegisterPassable。
若干标准库 API 获得了统一闭包重载:parallelize() 和 parallelize_over_rows()(在 std.algorithm.backend.cpu.parallelize 中)、bench.bencher()、DeviceContext.execution_time() 和 DeviceContext.enqueue_function()(GPU 入队路径,从之前的 enqueue_closure() 重命名)。
将 std.reflection 中的反射 API 统一到返回 Reflected[T] 句柄的 reflect[T]() 入口点背后。reflect() 通过预导入自动可用。句柄上的方法取代了 struct_field_* 系列自由函数(去掉了 struct_ 前缀——只有结构体才有字段)以及 get_type_name() / get_base_type_name() 自由函数:
struct Point:
var x: Int
var y: Float64
def main():
comptime r = reflect[Point]()
print(r.name()) # "Point"
print(r.field_count()) # 2
print(r.field_names()[0]) # x
comptime y_type = r.field_type["y"]() # Reflected[Float64]
print(y_type.name()) # "SIMD[DType.float64, 1]"
print(reflect[List[Int]]().base_name()) # "List"
var v: y_type.T = 3.14
旧版自由函数和 ReflectedType[T] 包装器现已 @deprecated;它们将在未来版本中移除。
align_down() 和 align_up() 现在接受泛型 SIMD[dtype, width] 整数值,取代先前的仅 UInt 重载。
扩展 FastDiv 和 mulhi() 以支持 64 位整数类型。
新增 Variadic.contains_value comptime 别名,用于在编译时检查可变参数序列是否包含特定值。