Mojo nightly
This version is still a work in progress.
Language enhancements
-
mojo now picks
Array(instead ofList) as the default type to construct from a list expression. E.g.,var x = [1, 2, 3]# type_of(x) = Array[Int, 3] -
whereclauses now accept an optional string-literal message, writtenwhere (condition, "message"). The message is included in the compiler diagnostic when the constraint fails, and is supported everywherewhereclauses are allowed: trailing function and struct constraints, struct conditional-conformance clauses, andalias/comptimedeclarations.def foo[sc: Int]() where (sc > 1, "scaling factor must be greater than 1"):...Calling
foo[0]()now reports the message in the note:note: constraint declared here evaluated to False, expected '(sc > Int(1))':scaling factor must be greater than 1The message must be a string literal; a non-literal message is reported as an error.
-
Support for
lambdaexpressions: anonymous, single-expression closures that desugar to a nesteddef. As in Python, the body is a single expression with noreturn; unlike Python, the arguments are parenthesized and typed like in adefsignature — for examplelambda (x: Int) {} -> Int: x + 1. The capture list{…}and return type may each be elided: an omitted capture list imm-captures the body's free variables (and is thin when there are none), and an omitted return type defaults toNone— so the barelambda: expris valid whenexprisNone-typed. These are fixed defaults, not inference (a non-Nonebody still needs an explicit-> T).A thin (capture-free)
lambdais a function value, exactly like adefreferenced by name. As such it:- binds to a
comptime; - passes as a
thinfunction-typed parameter; - decays to a
thinfunction pointer in runtime positions.
Referencing an enclosing function or struct parameter keeps it thin. Any other
lambdais a closure instance — a runtime value with no function type, so it does none of the above:- one that captures;
- one that writes an
{imm}/{mut}capture convention, even capturing nothing; - one with unbound parameters of its own (
lambda [N: Int](…)), bound at each call.
- binds to a
-
Mojo supports an (internal only for now) feature known as interior origins, which allows collections to protect from a common class of memory unsafety problems.
List, for example, now returns element references bound to an interior origin of the list instead of the whole-list origin, so an element reference is invalidated when the list is mutated (for example byappend()orpop()). Code that holds an element reference across such a mutation is now correctly rejected by the lifetime checker instead of silently dangling after a reallocation:var list = [1, 2, 3]ref elem = list[0]list.append(4) # may reallocate, invalidating `elem`print(elem) # error: use of invalidated interior reference -
Mojo now supports type inference from literal initializers:
var x: List[_] = [1, 2, 3]var y: List = [1.0, 2.0, 3.0] -
Mojo now supports
==and!=for type equality checks, and_type_is_eqhas been removed. -
Mojo now infers
TraitforTypeList.ofsuch thatcomptime TL = TypeList.of[Int, Bool]# works withoutcomptime TL = TypeList.of[Trait = AnyType, Int, Bool] -
Mojo now warns about redundant trait composition
# Warning: Redundant trait composition: 'Copyable' already implies 'AnyType'comptime T : AnyType & Copyable = xxx -
Keyword variadic arguments can now be forwarded to another function that takes keyword variadics, using Python style
**syntax:def takes_them(var **kwargs: Int): ...def pass_them(var **kwargs: Int):takes_them(**kwargs^) -
Dynamic function pointers with unbound type parameters can now be called directly. The compiler infers parameters from the call arguments and specializes the callee before the indirect call. This capability only works with a limited set of parameters - those which are specialized to a single value. This notably enables origin parameters on runtime function calls, which can also be implicit from variadics:
var fp1: def(*Int) thin -> Nonevar fp2: def[a: ImmOrigin](ref [a] x: Int) thin -> None...fp1(1, 2)fp2(42) -
Struct fields are no longer allowed to hide
UnsafeAnyOriginwithin a struct. For example, this is no longer accepted:struct Example:# error: cannot use UnsafeAnyOrigin in a struct field.var ptr: UnsafePointer[Int, MutUnsafeAnyOrigin]This is because Mojo doesn't know that uses of
Examplecontain anUnsafeAnyOriginand therefore doesn't do lifetime extension for values in its context. The typical solution for this is to add anOriginparameter but you can also useUntrackedOriginif you explicitly manage the lifetime of the underlying data:struct Example[origin: Origin]:var ptr: UnsafePointer[Int, Self.origin]# ORstruct Example:var ptr: UnsafePointer[Int, MutUntrackedOrigin]As a temporary workaround, you can decorate fields with
@__allow_legacy_any_origin_fieldsto ignore the compiler error, however this decorator is not stable and will eventually be removed. -
Method
selfparameters must now have typeSelf. Customselftypes are now rejected unless the method is annotated with the (temporary)@__allow_legacy_custom_self_typedecorator. Switch to awhereclause instead.struct Foo[T: AnyType]:# ERROR:def foo(self: Foo[Int]):...# Migrate tostruct Foo[T: AnyType]:def foo(self) where Self.T == Int:... -
Added support for checking variadic type-list operands with
conforms_to(). For example, a variadic parameter list can pass its type-list value directly:def copy_variadic_elements[*Ts: AnyType](*args: *Ts) where conforms_to(Ts.values, Copyable):passTo check several distinct standalone types against a trait, conjoin scalar checks, for example
conforms_to(T, Trait) and conforms_to(U, Trait). -
Mojo has improved its tracking of import locations and now shows where a package containing a diagnostic was first introduced into the program:
Included from /bug.mojo:2:Included from /foo/__init__.mojo:3:Included from /foo/nested_pkg/__init__.mojo:4:/foo/nested_pkg/my_module.mojo:1:5: note: candidate not viable: unexpected argumentdef bar(): pass^For precompiled packages (
.mojocfiles), locations inside the package are omitted. For brevity, the compiler also does not report wherestdpackages are pulled in, since they are implicitly imported into every module. -
immis now the preferred spelling for thereadargument and closure-capture convention.readstill works but will soon be deprecated. -
Parametric "generator" types can now be spelled with a dedicated keyword instead of having to use MLIR syntax directly. This keyword is subject to change in the future as we get experience with it. An example is:
def foo[type: __generator_type[size: Int] SIMD[DType.uint8, size*2]](): ....
Language changes
-
size_ofnow returns the allocation size: the store size rounded up to the type's alignment, which is the stride between adjacent elements of an array of that type. This changes the result only for types whose store size is not a multiple of their alignment (e.g., structs + whose@align(N)exceeds their natural alignment) and fixes memory corruption when such types were used inListand other collections whose growth copiescount * size_ofbytes. -
Mojo now rejects function overloads that differ only in argument convention (
immvsmut). -
Predefined and reserved words (for example
class,del,match,yield) can no longer be used as the name of a free function. Doing so now errors at the declaration instead of silently producing a function that could never be called. -
A bare
**kwargsis now an error; writevar **kwargs(a fixit inserts it), in function declarations and function types alike.varwas already the only supported convention — the sole exception to arguments defaulting toimm, applied silently before — so semantics are unchanged. -
User-written structs must now explicitly declare closure-trait conformance in their inheritance list to satisfy a
def(...) -> ...closure trait. Previously a struct with a compatible__call__was accepted implicitly (duck-typing). Declare the trait in the struct's inheritance list:def apply[F: def(Int) -> Int](f: F, x: Int) -> Int:return f(x)struct Double(def(Int) -> Int): # previously: `struct Double:`def __call__(self, x: Int) capturing -> Int:return x * 2_ = apply(Double(), 5)Conformance is checked at struct definition rather than deferred to the use site.
-
The import system has been overhauled to make name resolution explicit and consistent:
-
Import resolution now follows a consistent preference order within a directory: source packages, then precompiled
.mojocfiles, then source modules, then legacy precompiled.mojopkgfiles. Previously the order was unspecified. -
Relative imports must use
from(from . import foo); theimport .fooform is no longer accepted. -
Absolute imports
import a.b.cnow bind all ofa,a.b, anda.b.cinto the scope, where previously onlya.b.cwas made available. Two related bugs are fixed:import afollowed byimport a.bno longer errors with "invalid redefinition of 'a'", and function-scoped dotted imports (import a.binside a function body) now work. -
An imported package's submodules are now only accessible when the package's
__init__.mojore-exports them (for example, withfrom . import sub). An absolute import of the submodule (import pkg.submodule) always works, bypassing the__init__.mojo. -
Intra-package accesses without explicit
imports are deprecated and will be removed in a future release. A module must now explicitly import symbols defined elsewhere in its own package:# module2.mojo — uses foo() from __init__.mojo and module1.bar()from . import foofrom . import module1foo()module1.bar() -
Modules and packages can now be imported through regular (non-package) directories using the same path-like syntax, for example
import dir.nested_dir.module. An import statement that resolves to a directory cannot itself be used for scoped lookups (import dirthendir.nested_dir.module.foo()is an error). -
A standalone module can no longer import its own name (for example,
import utilinsideutil.mojo). Such an import could only resolve to the module itself, silently shadowing any same-named package on the search path. Modules inside packages are unaffected. -
Importing functions with the same name from different modules, combining them into one overload set, is now deprecated and emits a warning; a future release will reject the second import. Import the name from a single module instead.
-
Wildcard imports now resolve latest first, textually: declarations imported last shadow earlier ones, including those implicitly imported from
std.prelude. -
Error diagnostics on failed imports are now emitted per import site, instead of once per module.
-
-
The
@explicit_destroydecorator is no longer sufficient for astructtype to opt out ofDeinitableconformance. As before, all structs implicitly conform by default; to narrow or opt out, write a constrainedDeinitable where ...conformance instead —where Falsefor types that are never deletable, or a non-trivial condition:struct NeverDeletable(Deinitable where False):def destroy(deinit self):passstruct Container[T: AnyType](Deinitable where conforms_to(T, Deinitable)):var value: Self.TUsing
@explicit_destroywithout an error-string argument is now an error on bothstructandtraitdeclarations, since it has no effect; remove it.@explicit_destroy("custom error")can still be used to give users additional instruction when an instance cannot be deleted implicitly. -
The destructor dunder method should now be spelled
__deinit__, for naming parity with__init__. The old__del__spelling still works but now emits a deprecation warning with a fix-it to rename it:struct Example:def __deinit__(deinit self):pass -
whereclauses inside a parameter list (for example,[x: Int where x > 0]) are no longer supported, following a period of deprecation. Use a trailingwhereclause after the signature instead:# Old (no longer supported):# def foo[x: Int where x > 0]():# New:def foo[x: Int]() where x > 0:pass -
The compiler now rejects newlines in the middle of certain statements, where they were previously permitted:
- Between
def/struct/trait/comptimekeywords and the following identifier - Between the
asyncanddefkeywords on function definitions - Anywhere in the midst of an
importstatement, save for parenthesized import lists.
- Between
-
Struct types are now Movable by default. To opt-out of always-on movability, either explicitly specify a conditionally Movable conformance using
Movable where <cond>, or opt out of Movable conformance entirely usingMovable where False.
Library stabilizations
-
trait Deinitable -
trait Movable -
trait Copyable -
trait ImplicitlyCopyable -
List
def __init__(out self)def __init__(out self, *, capacity: Int)def __init__(out self, *, copy: Self) where conforms_to(Self.T, Copyable):def __init__(out self, *, length: Int, fill: Self.T) where conforms_to(Self.T, Copyable):def __del__(deinit self) where conforms_to(Self.T, Deinitable):def reserve(mut self, capacity: Int):def resize(mut self, length: Int, fill: Self.T) where conforms_to(Self.T, Copyable & Deinitable):def __getitem__[origin: Origin, //](ref[origin] self, slice: ContiguousSlice) -> Span[Self.T, origin_of(self)._get_owned_interior["element"]]:def __iadd__(mut self, var other: Self, /) where conforms_to(Self.T, Copyable):def extend(mut self, var other: Self):def __contains__[dtype: DType, //](self: Span[Scalar[dtype], _], value: Scalar[dtype]) -> Booldef __contains__(self, value: Self.T) -> Bool where conforms_to(Self.T, Equatable)def __getitem__(ref self, idx: Int, /) -> ref[_] Self.T:def __eq__(self, other: Self, /) -> Bool where conforms_to(Self.T, Equatable):
-
Bool
-
Span
def __init__(out self):def __init__(other: Span, out self: ImmSpan[other.T, other.origin]):
-
String
def __init__(out self, data: StringLiteral, /):def __init__(out self, *, from_utf8_lossy: Span[Byte, _]):def __eq__(self, rhs: String) -> Bool:def __eq__(self, other: StringSlice) -> Bool:def __ne__(self, other: StringSlice) -> Bool:def __getitem__(ref self, idx: Int, /) -> ref[self.origin, self.address_space] Self.T:
-
Optional
def __init__(out self):def __init__(out self, var value: Self.T) where conforms_to(Self.T, Movable):def __bool__(self) -> Bool:
-
Array
def __getitem_param__[idx: Int, /](ref self) -> ref[self] Self.T:def __getitem__(ref self, idx: Int, /) -> ref[self] Self.T:def unsafe_ptr[...](ref[origin, address_space] self) -> Pointer[...]:
-
ImmPointer
-
MutPointer
Library changes
-
The second parameter of
SIMDhas been renamed fromsizetolength, to match theSIMDLengthtype it is declared with and thelengthvocabulary the rest of the library uses for element counts:var v = SIMD[DType.float32, length=4](1.0, 2.0, 3.0, 4.0)print(v[0], Int(v.length))Positional uses such as
SIMD[DType.float32, 4]are unaffected. Reading the parameter asv.sizestill works but is deprecated and warns, so existing code keeps compiling while you migrate. Binding it by keyword asSIMD[dtype, size=4]is an error and must be updated.ComplexSIMD's matchingsizeparameter has been renamed tolengthas well, so the two types stay consistent. -
The
MutSpanandImmSpanaliases are now exported from the prelude, so they no longer need an explicit import fromstd.collections. This matches theMut/Immaliases forPointer, which the prelude already exported. -
StringSlicenow hasMutStringSliceandImmStringSlicealiases, matching theMut/Immaliases already provided forSpanandPointer. A function can use them to state the mutability it requires of a string argument without spelling out theoriginparameter. -
GPUInfo.vendorhas been removed. It duplicatedGPUInfo.api, which identifies the vendor precisely ("cuda","hip","metal", or a stdlib plugin's own API name) rather than collapsing every plugin accelerator into one enum value. Compareapiinstead:comptime use_apple_path = ctx.default_device_info.api == "metal"Vendoritself remains, as the classifier behindhas_amd_gpu_accelerator(),has_nvidia_gpu_accelerator()andhas_apple_gpu_accelerator(). -
ImplicitlyDestructiblehas been renamed toDeinitable, for consistency with thedeinitargument convention and the__deinit__spelling of the destructor. BothImplicitlyDestructibleand the intermediateImplicitlyDeletablespelling remain available as deprecated aliases. -
Spannow has a keyword-onlyaddress_spaceparameter (defaulting toAddressSpace.GENERIC), so a span can view memory in a non-default address space, such as GPU shared memory:var smem = stack_allocation[32, Float32, address_space = AddressSpace.SHARED]()var tile = Span[mut=True, Float32, MutUntrackedOrigin, address_space = AddressSpace.SHARED](unsafe_ptr=smem, length=32)Address-only operations (indexing, slicing,
unsafe_ptr(),as_imm(), and the SIMD search helpers) work in any address space and preserve it in their results. Element-copying operations (iteration,fill(),copy_from(), hashing, equality, and writing) remain restricted to the default address space, since copying a value into or out of another address space requires the element type to be trivially copyable. -
Any integer scalar can now be constructed from an
Intablevalue, not justInt. This makes taking a pointer's address as an unsigned integer work directly:var x = 42var p = Pointer(to=x)var addr = UInt(p) # previously required `UInt(Int(p))` -
Bencher.iter()now accepts a raising closure as a runtime argument, so a benchmark whose body raises can pass a closure with an explicit capture list instead of an@parameterclosure. Prefer the unified closure form over the deprecated@parameterone.def bench_add(mut b: Bencher) raises:var a = PythonObject(42)var c = PythonObject(10)@always_inlinedef call_fn() raises {var a, var c}:var r = a + ckeep(r)b.iter(call_fn) -
Arrayis no longerDefaultable. Previously it conformed toDefaultablebut attempting to actually default construct anArraywould fail to compile. -
Erroris nowImplicitlyCopyable, so re-raising a caught error no longer requires the transfer sigil:try:might_fail()except e:print("logging error:", e)raise e # previously an error: use `raise e^`A captured
StackTraceis now reference counted, so copying anErrorcosts a reference count increment rather than duplicating the trace.raise e^still works and avoids the copy. -
PythonObjectarithmetic, comparison, and membership operators now dispatch through CPython's abstract number, object, and sequence protocols (for examplePyNumber_Add,PyObject_RichCompare, andPySequence_Contains) instead of a Python-level attribute lookup followed by a bound-method call. Together with the non-mutating operators now borrowing their operand rather than taking it by value, this is roughly 12x faster on the interop hot path (a tighta + bora < bloop). It also follows standard Python operator semantics more closely, including reflected-operand fallback (__radd__,__rmul__, and so on) and the standard error messages for unsupported operations. An operation that no operand supports now raisesTypeError, where previously it could yield theNotImplementedobject as a value, and comparing mismatched types with==now returnsFalserather than a truthyNotImplemented. -
InlineArrayhas been renamed toArray, its first parameter fromElementTypetoT, and its second parameter fromsizetolength. A temporaryInlineArraycomptime alias exists for adoption, and.sizeremains as a deprecated alias for.length. Update explicitInlineArray[ElementType=..., size=N]usages toArray[T=..., length=N]. -
Many raw-pointer APIs across the standard library now use a safe
Pointerinstead of anUnsafePointer:List.unsafe_ptr(),InlineArray.unsafe_ptr(), andUnsafeUnion.unsafe_ptr().- The
unsafe_ptr()accessors ofSpan,StringSlice,String(plusString.unsafe_ptr_mut()),StringLiteral, andCStringSlice. Allocation.unsafe_ptr(),Allocation.unsafe_leak(), andOwnedPointer.unsafe_ptr().PythonObject.unsafe_get_as_pointer(),PythonObject.downcast_value_ptr(), andPythonObject.unchecked_downcast_value_ptr().- The AMD
sys.intrinsics.implicitarg_ptr()intrinsic. DevicePointer.unsafe_ptr()(std.gpu.host) and theunsafe_ptr()requirement of theDevicePointerLiketrait.- The
capture_sizesfield ofCompiledFunctionInfo(std.compile), now a safePointer[UInt64]. - The
Span(unsafe_ptr=..., length=...)constructor, matchingSpan's internal pointer field.
The two pointer types share the same layout and convert implicitly, so most code is unaffected. Code that called an unsafe-only pointer operation directly on the result should switch to the ungated
unsafe_*spelling, for exampleptr + ibecomesptr.unsafe_offset(i)andptr[i]becomesptr[unsafe_offset=i]. -
OwnedPointer.steal_data(),ArcPointer.steal_data(), andList.steal_data()have been renamed tounsafe_take_allocation()and now return an owningAllocationinstead of a raw pointer. The methods keep anunsafe_prefix because the elements are handed over still initialized: deallocating does not run their destructors. Recover the previous raw pointer withunsafe_leak().List.steal_data()andOwnedPointer.steal_data()remain as@deprecatedmethods, whileArcPointer.steal_data()is removed outright: the reconstructingArcPointer(unsafe_from_raw_pointer=...)constructor now takes a pointer to the control block (obtained fromunsafe_take_allocation().unsafe_leak()) and no longer accepts the payload pointer thatsteal_data()handed out. -
OwnedPointer.take()has been renamed toOwnedPointer.into_inner(). The old name remains as a@deprecatedmethod and will be removed in a future release. -
Variant.take[T]()andVariant.unsafe_take[T]()have been renamed toVariant.unwrap[T]()andVariant.unsafe_unwrap[T](). The old names remain as@deprecatedmethods and will be removed in a future release. -
The
as_immutable()method onUnsafePointerand theget_immutable()method onSpan,StringSlice, andUnsafePointerhave all been renamed to a singleas_imm()method, embracing the shorterimmspelling for a consistent immutability API. The old names remain as@deprecatedaliases and will be removed in a future release. -
Added
runtime.initialize_runtime(), which initializes the Mojo runtime when Mojo code built as a shared library (mojo build --emit shared-lib) is called from a non-Mojo host program such as C or C++. In that situation no Mojomain()function runs, so the runtime was never initialized and parallel or asynchronous APIs such asparallelize()crashed. Callinitialize_runtime()before using any runtime-dependent API; the call is idempotent, and a single call covers all threads in the process. See Call a Mojo shared library from C or C++ for details. -
Added
List.try_index(), which returns the index of a value in a list (if present) without raising, and is comptime-compatible. -
When an unhandled error propagates out of
mainand no stack trace was collected, Mojo now prints a hint to setMODULAR_DEBUG=stack-trace-on-errorto enable stack trace collection, rather than printing only the error message. -
OptionalandVariantnow accept element types that are notMovable. Their element types are now bounded byAnyType, withMovable,Copyable, and related conformances conditional on the element types. A non-Movablevalue can be stored in place with the new closure-basedinit_with=constructors andVariant.set()overload, which construct the value directly into storage (placement-new) rather than moving it.deinit_with()on both types also no longer requiresMovable, so element types that are neitherMovablenorDeinitableare fully usable:@fieldwise_initstruct Pinned(Movable where False):var value: Intdef make() -> Pinned:return Pinned(7)var opt = Optional[Pinned](call=make) # construct in placevar v = Variant[Pinned, Int](call=make)v.set(call=make) # replace in place -
Optionalno longer conforms toIterator; it is now anIterablecollection of 0 or 1 elements.for value in optandfor value in opt^are unchanged, but code that used anOptionaldirectly as an iterator (for example callingnext()on it) no longer compiles and should iterate theOptionalinstead. -
Various datatypes have adopted interior origins (described under language enhancements above), including
List,Deque,Variant,String,Dict,LinkedList,OwnedPointer, andHostBuffer. A reference or view into one of these containers now carries an interior origin, so one held across a mutation is rejected by the lifetime checker instead of silently dangling after a reallocation. For example,HostBuffer.as_span()now returns aSpanbound to an interior origin of the buffer instead of the whole-buffer origin:var buf = ctx.enqueue_create_host_buffer[DType.float32](4)var s = buf.as_span()buf[0] = 1.0 # mutates the buffer, invalidating `s`print(s[0]) # error: use of invalidated interior reference -
BitSetgainedtest_range[bit_value: Bool, *, lo: Int, hi: Int], which efficiently tests that a bit range holds an expected value, and resizing constructors: theresized_from:keyword constructor zero-extends from a smaller set (and debug-asserts no set bits are dropped when shrinking), while a companion overload taking atruncate_set_bits: ()keyword argument truncates instead. -
Added
Tuple.consume_elements, which moves each element out of a tuple into a caller-provided closure one at a time. Destructuring such asa, b = t^copies each element, so it cannot take apart a tuple whose elements areMovablebut notImplicitlyCopyable;consume_elementstransfers ownership instead, mirroringVariadicPack.consume_elements.var t = ([1, 2, 3], [4, 5, 6]) # `List` is not `ImplicitlyCopyable`@parameterdef handler[idx: Int](var elt: t.element_types[idx]):print(len(elt))t^.consume_elements[handler]() -
More renames standardizing on
lengthoversize:TypeList.sizeis nowTypeList.length, andDeviceContextListis renamed toDeviceContextArraywith itssizeparameter nowlength(update explicitDeviceContextList[size=N]toDeviceContextArray[length=N]). Similarly,List.resize()andList.shrink()now takenew_lengthinstead ofnew_size, and thevalueargument ofList.resize()is renamed tofill, matchingList's constructor. The old names remain as deprecated aliases where applicable. -
Span's pointer-and-length constructor argument is renamed fromptrtounsafe_ptr, to flag that this construction path is memory-unsafe: the caller must ensure the pointer addresses at leastlengthvalid elements. UpdateSpan(ptr=..., length=...)toSpan(unsafe_ptr=..., length=...). -
List.capacityis now acapacity()method instead of a public field. This keeps the allocated capacity out of the stable public field surface, since it should only change indirectly through operations likeappend(). Replacemy_list.capacitywithmy_list.capacity(). -
Renamed
StaticConstantOrigintoImmStaticOrigin, to align with theImm-prefixed spelling used for the other immutable origins. The old name is still available as a deprecated alias and will be removed in a future release. -
range()has been reworked:- The
Int-based andScalar-based range types are unified into a singledtype-parameterized family, now thatIntisScalar[DType.int].range()withIntarguments behaves exactly as before. As part of this,range(...).__len__()always returnsInt, and asserts when an unsigned range's element count exceedsInt.MAXrather than silently clamping or wrapping; usebounds(), whose upper bound isNonein that case, for the size hint. - Floating-point iteration is now drift-free and reversible. Element
iis computed asfma(i, step, start), so forward and reverse iteration produce identical sequences across repeated calls and across any IEEE-754 platform at the same floating-point width. Previously a step that was not exactly representable, such as0.1, could drift and yield an extra forward element thatreversed()then dropped. reversed()now works on typed ranges such asreversed(range(Int16(1), 10, 2)). TheReversibleRangetrait gained an associatedReversedTypeiterator instead of hard-coding its__reversed__()return type, so every range flavor can conform and return its own reversed iterator.- Non-numeric element types (
Booland the narrow MX float formats) are now rejected at construction, and the one- and two-argument float ranges (range(Float64(4.5))andrange(Float64(0.5), Float64(3.0))) are compile errors instead of infinite loops; use the three-argument stepped form.
- The
-
repr()of a scalarSIMDvalue (size == 1) now prints using its type alias instead of the verboseSIMD[DType.<dtype>, 1](...)form when the dtype has one. For example,repr(UInt32(4))is nowUInt32(4)(previouslySIMD[DType.uint32, 1](4)), andrepr(List[UInt](1, 2))is nowList[SIMD[DType.uint, 1]]([UInt(1), UInt(2)]).size > 1values, and scalar dtypes without an alias (such asDType.bool), keep theSIMD[...]form. This only affectsrepr();String(...)/print(...)output is unchanged. -
Renamed the raw memory functions to make their unsafety explicit:
memmove,memset,memset_zero,memcmp,uninit_move_n,uninit_copy_n, anddestroy_nare nowunsafe_memmove,unsafe_memset,unsafe_memset_zero,unsafe_memcmp,unsafe_uninit_move_n,unsafe_uninit_copy_n, andunsafe_destroy_n. The old names are deprecated and will be removed in a future release. -
Added
Dict.insert(key, value)andDict.clear_with(destroy_func), with mirroringSet.insert(element)andSet.clear_with(destroy_func), so aDictorSetwhose key, value, or element type is notDeinitablecan be populated and cleared. Unlikedict[key] = value,insertdoes not destroy a displaced entry: it moves it out and returns it as anOptionalfor the caller to destroy.clear_withhands each entry todestroy_funcand retains capacity:var d = Dict[Int, Int]()var displaced = d.insert(1, 10) # None — key 1 was absentdisplaced = d.insert(1, 20) # the displaced (1, 10) entry -
Dict.fromkeys(keys, value)has been generalized from taking aListto accepting any iterable of keys. Both forms require the key and value types to beDeinitable. -
By-reference
Dictiteration (for entry in dict,keys(),values(),items(), andreversed()) no longer requires the key and value types to beDeinitable. These iterators only borrow references and never destroy an entry, so they now work on aDictwhose key or value type is notDeinitable. Consuming iteration (for entry in dict^andtake_items()) still requiresDeinitable, since it drops the entries it does not yield. -
Spanhas moved fromstd.memory.spantostd.collections.span. -
The container backing variadic
**kwargshas been renamed fromOwnedKwargsDicttoStringDict.StringDictno longer requires its value typeVto beDeinitable. A keyword dictionary whose values are linear (non-Deinitable) is itself linear and must be torn down explicitly with the newdeinit_with(deinit_func), which hands each key and value todeinit_func. It also gainedinsert(key, value)(returns the displaced entry as anOptional[DictEntry]without destroying it) andpopitem()(moves out and returns a whole entry), mirroringDict. Operations that destroy a displaced value in place —kwargs[key] = valueand the two-argumentpop(key, default)— still requireVto beDeinitable; useinsert,popitem, or the single-argumentpop(key)for linear values. -
Coordnow conforms toDevicePassable, so aCoordembedded in aDevicePassabletype (such as aTileTensor'sLayout) is encoded to the device throughCoord._to_device_typeinstead of a raw field bit-copy, the same wayIndexListalready was. -
Added
copy_to_numpy_arrayandfrom_numpy_arrayto the newpython.numpymodule for moving flat numeric data between MojoSpan/Listand NumPy arrays without hand-writtenctypesplumbing:from std.python.numpy import from_numpy_array, copy_to_numpy_arrayvar values: List[Float64] = [1.0, 2.0, 3.0]var array = copy_to_numpy_array(values) # NumPy array (copies)var span = from_numpy_array[DType.float64](array) # borrow array as a SpanBoth support the fixed-width numeric dtypes.
copy_to_numpy_arraycopies its input into a new, independent array;from_numpy_arrayborrows the array's buffer zero-copy. -
Intis now an alias forScalar[DType.int]and integer literals materialize to thisScalartype. Because of this some conversions have become more strict.A new
SIMDLengthtype has been added for the width ofSIMDitself and must be used when inferring a parameter based on a SIMD argument like so:def frob[w: SIMDLength](v: SIMD[DType.int, w]): ...Alternatively the width can be unbound if you simply want to be parametric over any
SIMDtype:def frob(v: SIMD[DType.int, _])The new
Intshould still be used in all other situations.This type was briefly named
SIMDSizeearlier in this nightly cycle;SIMDSizeremains as a deprecated alias forSIMDLength. -
chdirhas been added to thestd.osmodule and anfchdirmethod has been added toio.FileDescriptor. These are wrappers for the corresponding POSIX functions. -
TypeList.all_conforms_to()is now implemented in terms ofconforms_to(), which supports parameter-list operands likeTs.values. As a result,all_conforms_to()constraints preserve the same proof structure as directconforms_to(Ts.values, Trait)constraints, so the compiler can use them in conditional conformance implication checks and type refinement.This means conditional conformances can rely on trait hierarchy relationships for an entire type parameter pack. Given a trait
JsonSerializablethat inherits fromSerializable, a conditionally conforming type previously had to repeat the inherited condition; now the derived condition alone is enough for the compiler to prove the inherited conformance:struct Packet[*Ts: Movable](- Serializable where Ts.all_conforms_to[Serializable](),JsonSerializable where Ts.all_conforms_to[JsonSerializable](),Movable,):passThe same constraints now refine each element of a variadic type parameter pack inside
where,comptime assert, andcomptime ifcontexts:def write_all[*Ts: Movable](mut writer: Some[Writer], *args: *Ts):comptime if Ts.all_conforms_to[Writable]():comptime for i in range(args.__len__()):args[i].write_to(writer) -
is_trivially_destructible()has been renamed tois_trivially_deletable(). It now accepts any type (T: AnyType) instead of requiringT: Deinitable, returningFalsefor non-Deinitable(linear) types. -
List.insert()andLinkedList.insert()no longer normalize negative indices. Mojo collections are moving away from negative indexing, so the valid index range is now[0, len(self)]; a negative index is out of bounds and aborts (checked when asserts are enabled). -
The
Reflected.field_type[name]reflection member has been renamed toReflected.field[name], because it returns a chainableReflectedhandle for the named field rather than the field's bare type, so the old name was not accurate. Retrieve the field's type from the handle's.Tmember, as inreflect[T].field["x"].T. A by-index dual,reflect[T].field_at[idx], has also been added so a field's concrete type can be recovered while iterating fields by index (where the name is not available as a literal):comptime y_type = reflect[Point].field_at[1]var v: y_type.T = 3.14 # y_type.T is the concrete field type -
Array[T](the type formerly known asInlineArray[T]) no longer conforms toImplicitlyCopyable, since it is not inherently cheap to copy. It continues to conform toCopyable. -
Several collection types now conditionally conform to
Deinitable, conforming only when their element type does. This lets a collection hold non-Deinitableelements at all (previously such a collection failed to compile); a collection of non-deletable elements is itself linear and must be drained explicitly with the newdeinit_with()method, which calls a closure on each element:collection^.deinit_with(my_destroy_closure)For
Deinitableelement types — the common case — all of this is transparent, but generic code that takes one of these collections by value may now need& Deinitableadded to its element bound so the collection can be dropped:def foo[T: Movable & Deinitable, //](var arr: InlineArray[T, 3]):passAffected types, and the operations that still require
Deinitableelements:InlineArray: no remaining restrictions.Deque: element-destroying operations (append,appendleft,extend,extendleft,insert,clear,remove, and so on) and consuming iteration (for x in deque^).Dict: element-destroying and key/value-copying operations (__setitem__,setdefault,fromkeys,update,__or__,__ior__,pop,clear) and consuming iteration, so aDictwith linear keys or values can currently be constructed and torn down but not populated or mutated.LinkedList: onlyclearand consuming iteration, so aLinkedListwith linear elements can be populated (append,prepend,insert,extend) and torn down.LinkedList.insert()also no longer raises on an out-of-range index; likeList.insert(), it now aborts (checked when asserts are enabled).Tuple: a tuple with a linear element must be torn down withdeinit_with()or fully consumed withconsume_elements(). Generic code that stores aTuple[*Ts]with an unbounded pack may need& Deinitableon the pack bound.Set: the element bound loosened fromKeyElement & Deinitableto justKeyElement; element-mutating operations (add,remove,discard,clear) and consuming iteration still require deletable elements, so aSetwith linear elements can be constructed and torn down but not populated.OwnedPointer[T]: conforms only whenTdoes; a linearOwnedPointermust be consumed explicitly withinto_inner()(for aMovableT) orunsafe_take_allocation()rather than dropped implicitly.
Consuming iteration is conditional through the
IterableOwnedconformance; generic code bounded onIterableOwnednow rejects a non-conforming element type at the bound rather than failing later inside__iter__(). -
InlineArray's element type bound loosened fromMovabletoAnyType, so anInlineArraycan now hold a non-Movableelement type. TheMovableconformance is now conditional on the element: move construction (including list-literal construction such as[a, b, c]) requires aMovableelement, while indexing, by-reference iteration, and destruction do not. Code that usesMovableelement types is unaffected, since aMovableelement still yields a movable array. -
Optionalgainedinto_inner(), the owning dual oftake(). Both return the contained value and abort when theOptionalis empty, butinto_inner()consumes theOptional(deinit self) instead of leaving it empty, so it does not have to write back the empty state.take()is unchanged. -
Optionalgaineddeinit_assert_empty(), which destroys an empty linearOptionalwithout a caller-provided deinitializer, aborting in safe-assert builds if it is non-empty.Optional.map()andOptional.and_then()also now work when the element type is linear (notDeinitable): they move the contained value out and destroy the emptiedOptionalexplicitly, so a linear value can be transformed and handed back to the caller. -
It is now possible to iterate over owned elements in
List,Dict,InlineArray,LinkedList, andSetwhen the element type is notCopyable: theIterableOwnedconformance on these collections now requires onlyMovable & Deinitable, droppingCopyable.def iterate[T: Movable](var list: List[T]):# Consume elementsfor var x in list^:pass -
The implicit conversion constructors that cast an
UnsafePointertoMutUnsafeAnyOriginorImmUnsafeAnyOriginare now deprecated and emit a deprecation warning when used.UnsafeAnyOriginis an unsafe escape hatch that silently extends unrelated lifetimes and disables exclusivity checking, so it should never be applied implicitly. Prefer keeping a concrete origin; if you must discard it, make the cast explicit with theas_unsafe_any_origin()method. -
Removed the implicit constructors that converted an
UnsafePointerinto anOptional[UnsafePointer[..., UnsafeAnyOrigin]]. Constructing anOptional[UnsafePointer]now preserves the pointer's real origin instead of silently widening it toUnsafeAnyOrigin. Two call-site updates may be needed:-
Passing a concrete pointer where the parameter's origin is a genuinely fixed
MutAnyOrigin/ImmutAnyOrigin(typically C-FFI signatures) now requires an explicitas_unsafe_any_origin(). -
Because origins are now preserved, exclusivity checking applies to
memcpy()(and similar) calls whosedestandsrcderive from the same buffer. An intra-buffer copy that previously compiled now errors with "argument of 'memcpy' call allows writing a memory location previously writable through another aliased argument". Opt out by making one argument an unsafe any-origin (the non-overlap ofdestandsrcis already amemcpy()precondition):memcpy(dest=buf + dst_off,src=(buf + src_off).as_unsafe_any_origin(),count=n,)
-
-
coordis now a comptime expression, andcoord[DType]()has been renamed todyn_coord[DType](). Now one can just writevar my_coord = coord[1, 2, 3]to create aCoord[ComptimeInt[1], ComptimeInt[2], ComptimeInt[3]]. -
Removed
trait_downcast_var(). Improvements to type refinement based onwhere conforms_to(..)andcomptime assert conforms_to(..)make explicit value trait downcasting no longer necessary. -
The
ConditionalTypetype function instd.utils.type_functionsis now deprecated. Use the equivalent ternary expressionT if cond else Uinstead:# Deprecated:comptime Storage = ConditionalType[If=cond, Then=Int, Else=NoneType]# Use instead:comptime Storage = Int if cond else NoneType -
Added
raise_python_exception()tostd.python.bindings, which translates a MojoErrorinto a Python exception viaPyErr_SetStringand returns a nullPyObjectPtr. -
The Python binding APIs now use safe pointers: the
PyCFunctionFastcalling convention used byPythonModuleBuilder.def_py_c_function()forMETH_FASTCALLcallbacks declares its argument array as aPointer[PyObjectPtr, MutUntrackedOrigin], typed-self methods registered throughPythonTypeBuilder.def_method()declare their self parameter as aPointer[Self](for example,self_ptr: Pointer[mut=True, Self]), and the extension argument helperscheck_and_get_arg()andcheck_and_get_or_convert_arg()return a safePointer. The pointer types share the same layout, so the C ABI and behavior are unchanged; update the spellings in signatures and read borrowed arguments withargs[unsafe_offset=i]. -
Iterating over a
String,StringSlice, orStringLiteralnow yields grapheme clusters by default. Their__iter__()and__reversed__()methods return aGraphemeSliceIter, sofor c in my_string:produces what a user perceives as a single "character" on screen. The lower-level views remain available when you want them:codepoints()orcodepoint_slices()for Unicode scalars, andbytes()for raw UTF-8 bytes. -
The
Equatabletrait now allows for positional-only implementations, and arguments on implementers no longer need to match the trait exactly. -
PointerandUnsafePointerhave had theirtypeparameter renamed toT. -
The
UnsafePointerpointee-lifecycle methods are deprecated in favor of unified replacements that work on anyPointer, so callers no longer need to wrap safe pointers inMutUnsafePointer:init_pointee_move()andinit_pointee_copy()becomeunsafe_write(): pass the value by move (ptr.unsafe_write(value^)) or as thecopykeyword argument (ptr.unsafe_write(copy=value)).destroy_pointee()anddestroy_pointee_with()becomeunsafe_deinit_pointee(): call it with no arguments to destroy anDeinitablepointee, or pass a deinitializing closure to destroy a non-Deinitablepointee in place.init_pointee_move_from()becomesunsafe_write_move_from(src), which moves the value out of a source pointer into the uninitialized memoryselfpoints to (leaving the source uninitialized).
-
Pointergained explicitunsafe_-prefixed methods for operations that are individually unsafe — unchecked bounds, aliasing casts, moving or overwriting memory — rather than requiring the whole pointer to be typed unsafe:unsafe_offset(),unsafe_load(),unsafe_store(),unsafe_strided_load(),unsafe_strided_store(),unsafe_gather(),unsafe_scatter(),unsafe_as_noalias(),unsafe_address_space_cast(), andunsafe_take_pointee(). These methods work on anyPointer. The previous unprefixed names still work, but are now hidden from the generated docs and remain gated behind an unsafe pointer type; prefer theunsafe_-prefixed names going forward. Each method's docstring documents the exactSafety:requirements the caller must uphold. -
The unprefixed pointer methods that the
unsafe_-prefixed names above replace —__getitem__(),__add__(),__sub__(),__iadd__(),__isub__(),load(),store(),strided_load(),strided_store(),gather(),scatter(),bitcast(),address_space_cast(),take_pointee(), andfree()— now emit a deprecation warning when called. -
The pre-unification pointer aliases
UnsafePointer,MutUnsafePointer,ImmUnsafePointer,ImmutUnsafePointer, andOptionalUnsafePointerare now deprecated in favor ofPointer,MutPointer,ImmPointer, andOptionalPointer. The two pointer types were unified some time ago; the old names only existed for source compatibility with code written before that unification, and now emit a deprecation warning when used. Update type annotations and constructor calls to use thePointerfamily instead:# Deprecated:var ptr: UnsafePointer[Int, MutUntrackedOrigin]# Use instead:var ptr: Pointer[Int, MutUntrackedOrigin] -
Pointernow supports subtracting two pointers to compute the signed distance between them in elements of the pointee type, via the newoffset_from()method (analogous to Rust'soffset_from). The-operator does the same. Unlike the other pointer-arithmetic operators, which produce a new pointer and stay gated behind an unsafe pointer type, subtracting two pointers returns anIntdistance and is available on safe pointers too:var ptr = alloc[Int32](4)var end = ptr + 3print(end - ptr) # => 3print(ptr.offset_from(end)) # => -3ptr.free() -
OwnedDLHandle.get_functionnow returns a callable that keeps the owning handle alive while it runs, fixing a crash where the library could bedlclosed between symbol lookup and the call. Its parameter is now the return type instead of the full function-pointer type, and it raises if the symbol is missing (previously it aborted the process):# Before:var sqrt = lib.get_function[def(Float64) abi("C") -> Float64]("sqrt")# After:var sqrt = lib.get_function[Float64]("sqrt") -
OwnedDLHandle.get_functionandOwnedDLHandle.callnow forward arguments using the C ABI rather than the Mojo calling convention, so structs can be passed and returned by value. Multi-field struct arguments are no longer rejected at compile time.
Tooling changes
-
Crash reporting now defaults to the
telemetry.enabledsetting, so the two are enabled or disabled together unless overridden. Settingcrash_reporting.enabled(or theMODULAR_CRASH_REPORTING_ENABLEDenvironment variable) explicitly still takes precedence. Previously crash reporting was disabled by default in one initialization path and enabled by default in production builds in another. -
The
program.crash_reporting_enabled_invocationtelemetry event has been renamed toprogram.initialized. It is emitted once per process whenever telemetry is enabled and carries acrash_reporting.enabledattribute recording whether crash reporting was on for that session. -
Added a
--lld-pathCLI flag. This overrides the LLD path that Mojo uses. -
mojo-lsp-serverno longer parses or type-checks code blocks inside docstrings by default. This checking rests on unstable foundations in the LSP server and was prone to failing, producing false-positive diagnostics unrelated to the code being edited, for little value in return. Pass-check-docstringswhen launchingmojo-lsp-serverfrom the command line to re-enable the previous behavior. We plan to make this checking more robust and re-enable it by default over time. -
Added a
--fp-modeCLI flag that controls floating-point behavior as a comma-separated list of items. The only supported feature now iscontract, one offast(default) oroff.contract=fastis like Clang's-ffp-contract=fast:a + b*ccan fuse into a fused multiply-add across statements and breaking strict IEEE compliance;contract=offdisables contraction for stricter floating-point semantics. The samecontract=fast|offitem is also accepted in theemission_optionof akgen.compile_offloadoperation, to control contraction of an individual offload kernel. -
Failed imports are no longer cached and may be retried, e.g., in the REPL.
GPU programming
-
IntandUIntno longer conform toDevicePassableand can no longer be passed as arguments to GPU kernels (viaDeviceContext.enqueue_functionorcompile_function). They are platform-sized index types whose bit width depends on the host, so passing them to an accelerator miscompiles when the host and device disagree on the width (for example a 64-bit host driving a 32-bit GPU index domain). Use a fixed-width type —Int32,Int64,UInt32, orUInt64— for kernel scalar arguments and parameters, and convert back withInt(...)inside the kernel body if you need a platformIntthere. A kernel that still takes a bareInt/UIntargument now fails to compile with: "Int and UInt are not passable to device kernels; use a fixed-width type such as Int32 or Int64 instead". -
Added programmatic Metal GPU frame capture in
std.gpu.host:_start_metal_trace_capture(ctx, path)and_end_metal_trace_capture(ctx)bracket GPU work and write a.gputracefile for offline replay (requiresMTL_CAPTURE_ENABLED=1). A_set_metal_gpu_print_enabled(ctx, enabled)toggle and theMODULAR_DISABLE_METAL_GPU_PRINTenvironment variable disable Metalos_logGPU print; print is also suppressed during a capture, which otherwise cannot be replayed. -
A bare
--target-acceleratorarchitecture (for examplegfx950orsm_90) is now handled identically to its vendor-prefixed form (amdgpu:gfx950,nvidia:sm_90). Previouslyhas_amd_gpu_accelerator(),has_nvidia_gpu_accelerator(), andhas_apple_gpu_accelerator()only recognized the vendor-prefixed spelling, so code that specialized on them (such as warp-tiling parameters) could silently take the wrong path and fail a downstreamcomptimeconstraint.amd:<arch>is also now accepted as an alias foramdgpu:<arch>, mirroring the existingnvidia:<arch>prefix. -
The GPU
Vendortype can now be imported fromstd.sys(from std.sys import Vendor). It remains importable fromstd.gpu.host.infofor backward compatibility. -
DeviceContext.load_functionnow keys its runtime cache on the requested entry-point name as well as the blob. Loading two different entry points (for examplekernel_aandkernel_b) from a single PTX/cubin blob no longer collides — previously the second load silently returned the function resolved by the first. The cache also no longer keys on the entire blob when no module name is supplied: it keys on a short hash of the blob instead, so each call avoids copying, hashing, and byte-comparing the whole blob (and retaining a duplicate of it). The win scales with blob size and matters most for large multi-entry blobs loaded on the per-execution path. -
The
DeviceStreamtype is now included in the API reference documentation. Returned byDeviceContext.create_stream()andDeviceContext.create_external_stream(), it provides methods for synchronizing and sequencing asynchronous GPU work (for example,synchronize(),record_event(), andenqueue_wait_for()). The type was already public but was previously hidden from the generated docs. -
Added an 8x8
simdgroup_matrixmatrix multiply-accumulate primitive (_mma_apple_8x8()) withapple_mma_load_8x8()/apple_mma_store_8x8()fragment helpers for Apple Silicon GPUs instd.gpu.compute.arch. Unlike the 16x16 path (Apple M5 only), the 8x8 primitive is available on all Apple GPU generations (M1-M5). It acceptsFloat16,BFloat16, andFloat32inputs with aFloat32accumulator. -
Atomic.compare_exchange()now accepts aweakparameter, and requiresweak=Trueto compile on Apple GPU targets: AIR exposes no strong compare-exchange primitive, so Metal only lowers theweakform. This is safe for the common case of a CAS-retry loop, since a spurious failure just costs one extra iteration. Previously any use ofcompare_exchange(), including helpers built on it like atomic scatter-reduce, failed to compile on Metal. -
Apple M5
simdgroup_matrixMMA now accepts FP8 (float8_e4m3fn,float8_e5m2) inputs with an F32 accumulator, alongside the existing F16/BF16/F32 and 8-bit integer types. -
Added
warp.match_any()andwarp.match_all():match_any()returns, for each warp lane, the mask of lanes whose value has the same bits, andmatch_all()returns the warp's active-lane mask if every lane holds the same bits and 0 otherwise. They use NVIDIA'smatch.any.syncandmatch.all.syncinstructions, areadfirstlaneballot fold on AMD, and a shuffle-based emulation on Apple Silicon GPUs. -
warp.vote()now works on Apple Silicon GPUs. Metal's AIR backend exposes no usable ballot intrinsic, so it emulates the ballot with an XOR-butterfly OR-reduction oversimd_shuffle_xor, returning a 32-bit mask (or aDType.uint64mask whose upper 32 bits are always zero); NVIDIA and AMD are unchanged. -
DeviceGraphBuilder.collect_dependenciesnow accepts an optionaldependenciesargument. The named predecessor handles are injected as ambient predecessors of every node theworkclosure adds, so the scope's nodes run after those predecessors without the closure threading the handles through to eachadd_*call. With the default (empty)dependenciesthe behavior is unchanged. Whenworkadds no nodes, the returned join node falls back to depending ondependenciesso it still chains correctly.var producers = builder.collect_dependencies(add_producers)# Every node added by `add_consumers` depends on `producers`:var consumers = builder.collect_dependencies(add_consumers, dependencies=[producers]) -
DeviceGraphBuilder.add_functionnow covers every liveDeviceContext.enqueue_functionform, so any kernel launchable on a device context can also be recorded as a graph node:-
Added an overload that takes the kernel as a compile-time parameter and compiles it automatically, so callers no longer need a separate
DeviceContext.compile_functionstep:def build(mut builder: DeviceGraphBuilder) raises {read}:_ = builder.add_function[kernel](42, grid_dim=1, block_dim=1, dependencies=[]) -
Added overloads accepting a
DeviceExternalFunctionloaded from PTX/SASS viaDeviceContext.load_function(), and a capturing kernel as a compile-time parameter with runtime arguments. -
All
add_functionoverloads now accept alocationargument so wrappers can attribute launch errors to their callers, and the closure overload now accepts (and honors) afunc_attributeargument.
-
-
Some standard library APIs related to accelerator programming have moved to a new
maxMojo package, including:std.benchmark.Bench.bench_multicontext->max.benchmark.bench_multicontextstd.benchmark.Bencher.iter_custom(DeviceContext)->max.benchmark.bencher_iter_customstd.gpu.compute->max.gpu.computestd.gpu.host->max.gpu.hoststd.gpu.memory->max.gpu.memory
-
AddressSpaceis now target-extensible rather than a fixed, portable enum. The built-in GPU spaces (GENERIC,GLOBAL,SHARED,CONSTANT,LOCAL,SHARED_CLUSTER,BUFFER_RESOURCE) are unchanged, but accessing any other name — for example an accelerator-specificAddressSpace.SCRATCHPAD— now resolves through the active hardware backend instead of being a hard-coded compile error. The set of valid address-space names is the union of the built-in GPU spaces and whatever the active backend defines, so accelerator backends can provide their own named spaces (with their own values) only where they exist. A name that no backend defines remains a compile-time error. -
Added support for the Steam Deck's RDNA2 Van Gogh APU.
-
The
layoutpackage is now bundled with MAX instead of Mojo. -
The GPU device-side standard library now uses the unified safe
Pointertype throughoutstd.gpu(memory,compute,intrinsics,sync, andprimitives). Public signatures that previously took or returnedUnsafePointerare respelled to barePointer; sincePointerandUnsafePointershare representation and origin and decay implicitly, this is a type-identical change for callers. One visible difference:external_memory()now returns a safePointerinstead of anUnsafePointer. Code that performs raw pointer arithmetic on the result or builds aLayoutTensor/TileTensorfrom it can wrap it in an explicitly-typedUnsafePointer[...]at the call site.
Removed
-
Removed the deprecated
DeviceContext.compile_function_experimental()andDeviceContext.enqueue_function_experimental()methods, along with overloads that passed the kernel twice. UseDeviceContext.compile_function[func]()andDeviceContext.enqueue_function[func]()instead. -
Removed the
UInt-returning GPU indexing accessors (thread_idx_uint,block_idx_uint,block_dim_uint,grid_dim_uint,global_idx_uint,lane_id_uint,warp_id_uint). Use theInt-returningthread_idx,block_idx,block_dim,grid_dim,global_idx,lane_id, andwarp_idaccessors instead. -
Removed the
store_volatile()andload_volatile()intrinsics fromstd.gpu.intrinsics. UseUnsafePointer.store[volatile=True]()andUnsafePointer.load[volatile=True]()instead, which work across all supported GPU targets rather than NVIDIA only. -
Removed the deprecated
GPUAddressSpacealias forAddressSpace. UseAddressSpacedirectly. -
Removed the
DType.invalidsentinel alias. Code that used it to represent an absent or optional dtype should useOptional[DType]instead. Accordingly,DType._from_str()now returns anOptional[DType](Nonewhen the string does not name a dtype) rather thanDType.invalid. -
Removed positional indexing on
StringLiteral(literal[i]). It allowed out-of-bounds reads and was inconsistent with the[byte=],[codepoint=], and[grapheme=]indexing scheme used byStringandStringSlice. Use those keyword accessors instead (for example, on aStaticString). -
Removed the static
String.write()methods. Use the equivalentString()constructor instead, which accepts the sameWritablearguments (for example,String(a, b, sep=", ")instead ofString.write(a, b, sep=", ")). The memberwrite()methods that append to an existing string are unchanged.
Fixed
-
Targeting an MI250X now works. While normalizing the architecture name,
gfx90awas rewritten to the nonexistentgfx90aa, so both--target-acceleratorandGPUInfo.from_namereported every spelling of the target (gfx90a,mi250x,amdgpu:gfx90aandamd:gfx90a) as an unsupported architecture. -
Code completion now reports the correct completion kind for names bound by a
from module import namestatement that hasn't been resolved yet. Structs, traits, and functions imported this way previously completed with no kind at all. Additionally, a renamed binding (from module import name as other_name) no longer disappears from the completion list when another binding to the same declaration is in scope. -
debug_assertgenerates less code, so builds with-D ASSERT=allcompile faster. Calls with no message arguments no longer allocate a 2048-byte message buffer in the caller's frame, which previously grew with the number of asserts and could push GPU kernels past the stack frame limit. A no-message assert failure now reportsassertion failedinstead of an empty message. -
debug_asserthas dedicated overloads for the no-message case, which generate less code and so compile faster than passing an empty message list. -
Code folding in VSCode now works for Mojo files.
mojo-lsp-serverno longer advertises folding-range support, which only produced docstring ranges and caused VSCode to disable its built-in indentation-based folding — leaving functions, structs, and blocks unfoldable. Editors now fall back to indentation-based folding until the server returns structural folding ranges. -
Fixed
print()anddebug_assert()emitting garbled output on AMD GPUs when a printed string's byte length was an exact multiple of 8. The AMDGPUhostcallprintf interface reads each string up to its nul terminator, and the terminator was being dropped in that case, so the host read past the payload. -
base64.b16decodenow raises on invalid input instead of silently producing corrupt output. -
Closures mixing
*args, named keyword-only arguments, and**kwargsnow all work as values. A capturing closure taking**kwargsno longer fails to compile ("no matching method in call to '_insert'"), and a call may now combine a*unpack, literal keyword arguments, and a**splat, as in Python:f(*args, **kwargs^)forwards both packed variadics directly, andf(1, named=2, **kwargs^)binds the literal keyword to its own named parameter alongside the splat. The reverse splat order (f(**kwargs, *args)) is rejected, matching Python, as is combining a**splat with other keyword arguments bound for the same**kwargs. -
#6784, #6434 -
math.sqrtonFloat64now works on NVIDIA GPU. It lowers to the IEEE correctly-rounded hardware sqrt (sqrt.rn.f64) instead of being rejected at compile time. NVIDIA has no approximate f64 sqrt, so theFloat32fast path continues to usesqrt.approx.ftz.f32. -
#4473 - The
offsetparameter ofFileHandle.seek()(andNamedTemporaryFile.seek()) is now a signedIntinstead ofUInt64, so negative offsets relative toos.SEEK_CURoros.SEEK_ENDwork as the docstrings already showed. Previously a negative offset only compiled as a literal (via unsigned wrap-around) and could not be passed from a signed variable. -
#6755 - Volatile loads are no longer removed when their results are unused.
-
Type refinement from a
conforms_to()guard now applies inside the branches of a ternaryexp1 if cond else exp2used in acomptimecontext, matching the existingcomptime ifstatement behavior. For example,T.property if conforms_to(T, HasProperty) else 0now compiles; previously the true branch failed with'AnyType' value has no attribute 'property'becauseTwas not refined under the guard. -
A
comptimemember with a trailingwhereclause is now accepted as a witness for a conditional trait conformance when the conformance constraint implies the member's constraint, for example:trait StaticSize:comptime SIZE: Intstruct Foo[size: Int = -1](StaticSize where size >= 0):comptime SIZE: Int where Self.size >= 0 = Self.size -
The reflection-based default
Equatableimplementation no longer fails to compile for single-elementRegisterPassablestructs. Such a struct is flattened to its sole field's type, which previously caused the reflectionfield_refto produce an invalidkgen.struct.gep. -
A method whose return type references a constrained
comptimemember (one declared with a trailingwhereclause) is now accepted when the method's ownwhereclause discharges that member's constraint. -
A method whose return type is a generic struct instantiated with a parameter that only satisfies the struct's declared trait bound via the method's own
whereclause (rather than via the parameter's own declaration) is now accepted, instead of spuriously rejecting the returned value as a different, unconvertible type. -
A struct using
where Falseto opt out of a builtin trait's implicit synthesis (for example,Movable where False) no longer spuriously fails to compile when one of its fields also opts out of that same trait. -
CPython.PyCapsule_Newnow takes itsnameargument as aStaticStringinstead of an ownedString. CPython stores thenamepointer directly in the capsule rather than copying it, so an ownedStringargument left the capsule holding a dangling pointer once the temporary was destroyed. -
A failed import no longer poisons its name for the rest of the compilation. Previously, after something like
import pkg.utilfailed to resolve, a laterimport utilwould silently bind the cached failure even when a realutil.mojoexists on the search path, making the module unimportable with no diagnostic. -
#6485 -
Optional[T]andVariant[...]no longer corrupt data for payload types that include aBoolfield. The fix changes how unions are lowered to LLVM. -
Struct extensions are no longer imported onto structs which happen to share a name with their intended struct, when the extensions' intended struct is shadowed by another:
from pkg_a import * # defines a Foo and extensions on itfrom pkg_b import Foo # defines another Foo and extensions on itPreviously in the above example, the extensions defined by
pkg_awould be imported and callable on the unrelatedFoostruct imported frompkg_b. -
Importing a package whose name is a prefix of another package when split by dots no longer works:
# Used to import e.g., package_with.dots if it presented as a package:# package_with.dots/# └── __init__.mojoimport package_with # now errors -
Importing escaped-identifier packages & modules whose names contain dots now works reliably.
from `package.with.dots`.`module.with.dots` import foomojo docand file-in-package builds also now use the whole dotted name for such packages, rather than truncating it at the first dot. -
Invalid SIMD vector lengths are now rejected during code generation.