Skip to content

Releases: suidvandiewereld/Mettle

Mettle v0.17.0

Choose a tag to compare

@suidvandiewereld suidvandiewereld released this 01 Sep 13:18

Mettle 0.17.0

I've always had Mettle focus on being a compiler-centric project. The problems I wanted were in the backend: register allocation, vectorizers, a linker of its own, and getting all of it to run with no LLVM underneath. The language on top was mostly whatever the compiler needed in order to have something to compile.

0.17.0 changes that. I spent the last 2 weeks building and further fleshing out the language itself, and here's what I'm sharing with you!

Breaking changes

Read this section before upgrading. Code written against 0.16.3 may not compile.

Plain enums are opaque. An enum was an integer to the type checker. It widened to int32 without a cast, two unrelated enums compared as equal, and a Color went where a Role was declared. Enums now take part in no implicit conversion, a written cast is the only way across, and a switch on an enum names its variants.

print_int is gone. 408 call sites migrated. Use interpolation: println("{value}").

Narrow integer arithmetic wraps at its declared width. Only a store into a narrow location used to truncate, so big + big > 0 answered yes for two int32 values whose sum is negative, and interpolation printed 4000000000 for an int32. Lowering now cuts +, -, * and << back to width. Programs relying on the old 64-bit intermediate will compute different answers, and the new ones are the documented answers.

std/conv and std/net report failure through Result and Option.

Was Now
str_find, str_find_byte answered -1 Option<int64>
str_split_once answered a 3-tuple Option<StrSplit>
str_to_i64 answered a second int Result<int64, string>, with str_to_i64_or
net_init, socket_tcp, socket_udp, sockaddr_in, sockaddr_in_any, send_all Result, carrying the platform error code

Linux links natively by default. The internal ELF linker writes the image; ld remains the fallback and --linker gcc still forces the driver.

Uninitialized aggregates start zeroed, as the docs always promised. A string local used to begin as a wild pointer.

! answers a bool. A for initializer gets its own scope, as does a shadowing var. Arrays decay to their address in every pointer position; passing a bare array used to pass its first 8 bytes and silently no-op fgets(buf, ...).

Language

  • char, a distinct one-byte type. s[i], for c in s, and "{c}" prints text.
  • String interpolation, "{expr}" for any expression.
  • Result<T, E> and Option<T> in std/core. Absence and failure are different questions and each gets its own answer.
  • Multidimensional arrays. int32[3][4] is three rows of four; dimensions read left to right, so grid[i][j] takes i from the first. Both dimensions bounds-check, at compile time for a constant index and under --safe for a computed one.
  • Recursive tagged enums, two types may point at each other.
  • Generic type-argument inference from what the arguments already say.
  • ++ and --, slices carrying their extent (T[], T[..]), heap arrays that know their length, function-pointer arrays, comptime generation from a constant table, and cross-module export const.

Targets and linking

Shared libraries on Linux. The ELF linker binds a .so, emits one, and publishes a program's symbols so a library can call back in. raylib is the case it was built for.

mettle --build app.mettle -o app -L/opt/acme/lib -lacme --rpath /opt/acme/lib

New: -l, -L, --rpath, --dynamic-linker, --export-dynamic, --shared, --soname.

Bare metal. Inline assembly, volatile that keeps its guarantee, @naked and @interrupt entries, cross-compilation via --target, a chosen link address via --image-base, flat images via --emit-flat, and 16-bit code generation.

Native ELF images with program headers, section headers and a symbol table, behind a format-neutral object reader serving both COFF and ELF.

Correctness

Roughly 100 miscompiles and internal compiler errors fixed. The recurring themes were narrow integer width, unsigned shift and divide read as signed, shadowing across locals, globals, parameters, match arms and for initializers, aggregate and tagged-enum copies, closure float ABI, and stack alignment at heap calls.

Three worth naming, all found late and all shipped in earlier builds:

  • Vectorized SiLU and exp reprocessed their own output. The kernels ran a final overlapping vector to cover the remainder, which is sound writing into a separate destination and wrong in place: overlapped elements got f(f(x)). Exactly 8 - (n % 8) elements wrong for every n > 8 not a multiple of 8, a 57% error on SiLU. SwiGLU is a transformer feed-forward block, so the LLM engine was corrupted for any hidden dimension not divisible by 8. Every vectorizer test used a round count, so the remainder path had never been executed by anything.
  • A parameter assigned once in the body aliased its source. The call that passes an argument writes the parameter before any instruction runs, and that write has no IR to count, so var t: int64 = b; b = 5; return t; returned 5 and gcd returned 0 under -s.
  • The ELF linker rejected clang's GOT load. Only the mov form of a GOTPCRELX relocation could be relaxed; clang emits add rdx, [rip+disp32], so the native link failed, fell back to ld, and still exited 0.

The test corpus grew from 799 to 943 .mettle cases. The suite now runs every program it previously only compiled, proves it ran every case, compiles the codegen corpus at -O as well, adds a parallel differential harness over the examples, sweeps every in-place SIMD kernel by length rather than by round counts, and gates cyclomatic complexity against a recorded budget.

Diagnostics and tooling

Diagnostics render as a framed source table with syntax colour, through a shared style layer the --explain report also draws through. Crash locations order totally, so a fault names the same statement on both platforms. Crash reporting is on by default at function granularity for about 8KB and no codegen cost, with --no-crash-report to opt out. --dump-ast prints the tree. Every documentation snippet is verified against the compiler, and the code reference is generated from the compiler's own help tables behind a drift gate.

Performance

Compile speed roughly doubled on both platforms, all of it from removing quadratics: a declaration index in place of one scan per name, label and branch targets answered from an index, hash keys computed once, word-at-a-time string comparison, and an arena allocator. Generated code gained a real memory model in the redundancy pass, alias analysis backed by an index, load elimination over the dominator tree, SLP pairing of adjacent float64 statements, dense switch dispatch through a table, and spill costs charged by loop depth.

Full Changelog: v0.16.3...v0.17.0

Mettle v0.16.3

Choose a tag to compare

@suidvandiewereld suidvandiewereld released this 19 Aug 03:45

Mettle v0.16.3

28 commits since v0.16.2. The language surface is unchanged, so existing code compiles as before. This release is backend work: the loop recognizers gained a shared foundation and a rot gate, the register-allocated backend picked up most of what it previously bailed on, the compile-time interpreter grew a full memory and call model, and three correctness fixes landed.

Correctness

Byte vectorizer sign extension (release-only miscompile). A byte load's lane value is the zero-extended byte whatever the element's declared signedness, because every scalar backend loads bytes with movzx. Honoring the declared signedness led the kernel to select vpmovsxbd for int8 arrays, so a release build of if ((int32)a[i] > 100) saw -106 where a debug build saw 150. Fixed at both decode sites and covered by a regression test.

mingw gcc 16 dropped emulated TLS. MTLC_THREAD_LOCAL now degrades to a plain static where the toolchain no longer supplies emutls, so the compiler builds on gcc 16.

Fingerprint snapshot lifetime. The snapshot owns its label strings, and stddef.h is included where size_t is used.

Loop recognizers

The recognizer pipeline was split into canonical form, a worklist driver, and a tail stage, which replaced repeated full sweeps. Supporting that:

  • One strength-reduction table is now shared between the x86 and ARM64 backends. The table is proven against its own test, and the duplicate copy of the math in the backend was deleted.
  • Loop recognizers share an affine model. The arithmetic was hardened against overflow, the facts are computed lazily, and the third recognizer sweep is gone.
  • Canonical form is checked structurally, so a loop that merely stopped changing no longer passes as canonical.
  • A claim baseline (tests/loop_claims.baseline) fails the build when a kernel a recognizer used to claim goes unclaimed. Silent kernel rot was previously invisible. Regenerate with tools/regen-loop-claims.ps1.
  • A loop fingerprint diagnostic and an e-class pilot landed. The pilot was trimmed to what it actually measured, and the measurement is recorded alongside it.

Register-allocated backend

Functions that previously fell back to the older path now go through the allocated backend:

  • Heap allocation, global aggregates reached by name, and the byte-offset fill walk are lowered natively.
  • Global dirtiness is tracked flow-sensitively, so a function that both writes a global and makes a call can still be allocated.
  • MIR carries strings as the aggregates they are, and the string value convention is honored through loads, stores, and call arguments.
  • Float stack parameters are homed, float stack arguments are passed, integer literals passed to float parameters are folded at the call site, and any scalar coerces into a float argument.
  • Indirect calls are typed through temps.
  • rotate_add is lowered natively, popcnt is lowered, and vectorized reductions, scalar-reading vloops, and four more audited kernels cross the kernel bridge.
  • Nonzero fill starts fold into the kernel base and count.

Compile-time interpreter

ir_interp now models the whole memory and call model, and string literals have real backing memory. This widens what --verify and --pgo can evaluate at compile time instead of bailing.

Full Changelog: v0.16.2...v0.16.3

Mettle v0.16.2

Choose a tag to compare

@suidvandiewereld suidvandiewereld released this 18 Aug 07:24

Mettle v0.16.2

Programs are about eight times smaller. A release build of examples/fib was
74,752 bytes and is now 9,728. The machine code Mettle generates is unchanged,
so benchmarks and compile times are the same. Windows 931/931, Linux 917/917.

Four things were padding every binary.

Fixed

  • 33 KB of zeros were written into every image. GCC with -fdata-sections
    gives each zero-initialized static its own .data section carrying real zero
    bytes, and the runtime has several large ones. An all-zero data section with
    no relocations is now carried as .bss, which occupies no file bytes.
  • The internal linker kept every section of every input object. It now
    collects the ones nothing reachable references, following relocations out
    from the entry point and from the sections that cannot be collected.
  • Unused library functions and their string literals could not be dropped.
    Each function now gets its own section (.text$name on COFF, .text.name on
    ELF), and each string literal gets its own on COFF. A program that never
    calls print_int no longer carries it or the strings it names.
  • Switch jump tables rooted the whole printf family. GCC puts them in a
    plain .rdata section that has to be kept, and their relocations point into
    vsnprintf, vsscanf and the float formatters. The runtime is now built
    with -fno-jump-tables, worth about 9 KB in every binary.

Sizes

Release builds on Windows: fib 9,728 bytes, base64_encode 7,680,
crc32 and word_count 7,168, float_sum and binary_search 6,144.
A debug build of fib is 8,192, down from 73,216. On Linux fib is 8,840.
A --safe build of word_count is 16,384, its 64 KB of checked-access
tables having moved to .bss.

New

METTLE_LINK_GC_REPORT=1 prints every section the link kept or collected,
with its object and size, to stderr. This is what found the jump tables.

Changed

  • --emit-obj objects carry one section per function, so anything that
    reads them with an external linker or an object tool sees the new names.
    Linux links already pass --gc-sections and collect them.
  • An undefined external that nothing reachable references is no longer
    recorded.
    A program that declares an extern it never calls does not import
    it.
  • Building Mettle from source now compiles the runtime with -fno-jump-tables.

Full Changelog: v0.16.1...v0.16.2

Full Changelog: v0.16.1...v0.16.2

Mettle v0.16.1

Choose a tag to compare

@suidvandiewereld suidvandiewereld released this 17 Aug 20:53

Mettle v0.16.1

Linux now runs the same test suite as Windows. Windows 931/931, Linux 917/917.
Running it there for the first time found four real ELF backend defects, fixed
here. The ten cases Linux skips are PE, COFF and Win32 surfaces with no
counterpart.

Fixed

  • Structs passed to C by value were wrong on Linux. The backend applied the
    Microsoft x64 rule everywhere. System V eightbyte classification is now
    implemented, both directions, including C calling an exported Mettle function.
  • The ELF link carried no on-demand runtime, so comparing two strings failed
    at ld on an undefined mettle_string_eq.
  • Profile reports named the wrong functions on Linux: the function id went in
    ECX with Win64 shadow space instead of EDI. new, free and realloc reached
    for the Win32 heap on every target.
  • A native AArch64 Linux host was treated as Windows: .obj outputs, the COFF
    scanner on ELF objects, and std/io selected over std/io.linux.
  • debug.o is bundled on Linux. --explain prints ASCII when redirected.

New

--debug-hooks works on Linux, over a FIFO named by METTLE_DBG_PIPE. A
debuggable build still links no libc. The editor extension remains Windows-only.

Changed (Linux)

  • --build with no -o writes foo, not foo.exe. Most likely to break a
    build script; pass -o foo.exe or use the new name.
  • A source with no extension is refused, rather than writing the product over
    the input.
  • Internal functions take local linkage in whole-program builds, so a program
    may define close, read, write, send, stat or abort. --emit-obj
    objects keep every symbol global.

Full Changelog: v0.16.0...v0.16.1

Mettle v0.16.0

Choose a tag to compare

@suidvandiewereld suidvandiewereld released this 17 Aug 02:42

Mettle v0.16.0

The next generation.

250 commits since v0.15.1. Mettle is now a single repository: the compiler, the backend, the linker, and the runtime build from one tree with no pinned dependency between them. This release adds a checked-access memory-safety mode, compile-time reflection, verified hot code swapping, and gives string a real library after fixing the reason it never had one.


Breaking changes

Narrowing conversions need an explicit cast

Widening still happens silently; narrowing reports M0119 and stops the build.

var n: int64 = 300
var a: int32 = n           // M0119
var b: int32 = (int32)n    // says the wrap is intended

This will touch existing code. The diagnostic names the file, the line, and the range the destination type holds.

print and println take string, not cstring

Call sites passing a cstring need print_cstr / println_cstr.

== and != on strings compare contents

They previously compared the 16-byte record as a scalar and were always false, including "ab" == "ab". Code that worked around this with streq still compiles; code that silently never matched now matches.

Other

  • A parameter can no longer be shadowed by a local of the same name.
  • Heap allocation is typed at the boundary. Allocators hand out rawptr, which converts to and from every pointer type, so var a: int32* = malloc(n) and free(a) need no cast.
  • The editor extensions moved to their own project.
  • libmtlc is no longer a pinned dependency. libmtlc.version is gone and the backend compiles from src/. It still ships as a release asset and the fetchers still work; what changed is that its API version (mtlc_version(), now libmtlc 0.2.0) is independent of the toolchain version. See Is libmtlc still a separate artifact?

Memory safety: --safe

--safe checks every access the compiler cannot prove in bounds, and reports what survived.

The work went into making the checks cheap: indices resolved against the heap, globals, and stack locals; an index recognized as coeff * counter + invariant + const; a single whole-loop check hoisted through bodies that branch and rejoin; a surviving check lowered to a comparison rather than a call. A checked function keeps its register allocation.

What it does not cover is written down in docs/memory-safety.md.


Hot code swapping

A function can be replaced in a running process, at a point the program names.

policy(5);                                  // 6
mettle_swap_stage(&policy, &policy_v2);
policy(5);                                  // still 6, staged but not applied
quiesce;
policy(5);                                  // 50

Staging records an intent; quiesce; is the only place it takes effect, so a replacement can never land mid-operation.

@swappable marks a function replaceable and keeps the call boundary a swap redirects: it implies @noinline, and @swappable @inline is refused, because an inlined body has no call to redirect.

The binding is a slot holding a function pointer, so applying a swap is one pointer-sized store. No code is modified, no page is made writable, and a thread already inside the old body finishes there.

Tool What it does
mettle swap-check old new runs the differential harness over two functions; refuses a changed signature
layoutof(T) layout digest, so a swap cannot silently reinterpret a value

A program with no quiesce; never links the swap runtime.


Compile-time reflection

comptime for walks a type's fields, Type and Field queries answer at compile time, and at module scope it generates declarations.

New in this release: fieldof(T, name) reaches a Field through a compile-time string, which is the one way into the field table a metaprogram can compose. Every other spelling needs the name written in source.

comptime for f in typeof(Packet).fields {
  total = total + (int32)fieldof(Packet, f.name).type.size;
}

Compile-time strings now compare, so a contract can span two declarations.

mettle expand prints the expansion as source, and mettle trace names the iteration behind each value instead of merging them:

21 |   total = total + f.offset ...   <- (field `kind`) total = 100; (field `seq`) total = 505

Strings

string had two representations at once: a value was a pointer to a {chars, length} record, a local was the record itself. Sites disagreed silently, and returning a string built in the callee wrote its fields through an uninitialized pointer. Nothing in the corpus returned a string, so nothing caught it.

string is now an aggregate and moves like the struct it is: copied whole, returned through a hidden pointer, stored inline.

That unblocked a real library. std/conv gains:

str_slice(s, start, len) the view at start, clamped to what s holds
str_starts_with / str_ends_with 1 or 0
str_eq_at(s, offset, needle) whether needle's bytes sit at offset
str_find / str_find_byte first index, or -1
str_contains 1 or 0
str_trim / str_trim_start / str_trim_end whitespace removed
str_split_once(s, sep) (head, tail, found)
str_to_i64(s) (value, ok); needs no terminator
i64_to_str(n, buf, buf_len) writes into buf, returns a view of it

Every returned string is a view into its input: nothing allocates, nothing copies.

var head: string = "";
var tail: string = "";
var found: int32 = 0;
(head, tail, found) = str_split_once("id=907", "=");

var value: int64 = 0;
var ok: int32 = 0;
(value, ok) = str_to_i64(tail);        // 907, no allocation anywhere

atoi, atol, cstr_len, and streq keep cstring as the C boundary.


Optimizer and vectorizer

The vectorizer reads shapes it previously declined:

  • an if that only selects a value, read as a value
  • counting under a predicate, without a branch
  • a comparison used as a value
  • a scan seeded from the first element
  • a global array's base hoisted above its loop
  • byte maps run in int32 lanes
  • a running maximum recognized as the operator it is

--explain reports what each loop and call became and why. Its fix suggestions are applied to a clone and re-checked before printing, and a guard that was discarding every proven inline fix is repaired.


Correctness

Each of these was a wrong answer, not a missed optimization:

  • == on strings was always false
  • returning a string built in the callee crashed, or returned garbage under --release
  • a predicated accumulate read the branch it had just retired
  • one arm of a select read the other arm's write
  • a call with many arguments built the wrong frame
  • three miscompiles around the integer range rule
  • defer did not run on break, continue, labeled jumps, or switch cases

Diagnostics

One mistake now produces one diagnostic, and the parser resyncs at block boundaries so a syntax error no longer cascades.

A GPU-only construct compiled for a CPU target used to report an internal compiler error naming an IR opcode number. It now names the construct and the flag:

error: 'tensor_mma' in function 'gemm_tile' runs on a GPU and has no CPU
translation. Compile the module that defines this kernel with --emit-ptx
(NVIDIA) or --emit-spirv (OpenCL), and keep it out of the host program

Runtime

Each optional component is independently excisable, and the absence is checked on every build rather than asserted. The swap and string runtimes are written in Mettle, compiled by the compiler that ships them.


Installer

The Windows installer is rebuilt in the language's own palette, with artwork generated from the brand mark at every display-scaling step, in light and dark. A guard fails the build if the version drawn on the banner is not the version being built. It reports the correct version and links to the right repository, both of which were stale.


Documentation

docs/ideology.md sets out the rules the language holds itself to and what each one costs. docs/memory-safety.md covers --safe, and the backend has its own reference under docs/libmtlc/.

The compiler-debugging pages documented four dump flags the compiler rejects; they now describe what exists. Eight working flags that appeared nowhere are documented, and every example presented as a complete program compiles.


Tests

The suite is at 927 checks, from 636 .mettle fixtures at v0.15.1 to 792.

Full Changelog: v0.15.1...v0.16.0

Mettle v0.15.1

Choose a tag to compare

@suidvandiewereld suidvandiewereld released this 31 Jul 15:58

One improvement to --report-occupancy, filed as feature request 9 by the
same inference engine after using the report in anger.

The occupancy report now shows whether a launch can fill the card

The per-SM residency ceiling says nothing about whether a kernel has enough
work to reach it. A 16-block launch on a 36-SM card is work-limited at any
residency, but the report read "50%, register-limited" and invited a tuning
detour that changed nothing.

When the SM count is known, the header now names it and every kernel line
carries the whole-card fill threshold, so the reader can put their grid size
next to it and see at a glance whether the ceiling is reachable:

Occupancy report (sm_121a; ... allocation unit 1; 36 SMs, local GPU):
  attention: 50 registers, block 32 (1 warps/block, 24 blocks)
    -> 24/48 resident warps (50%), register-limited; full card = 864 blocks (36 SMs x 24)

Sixteen attention heads against a 864-block threshold is the work-limited
verdict in one line.

The SM count comes from --sms=N, or from the local driver when the flag is
absent. The driver is loaded dynamically, so the compiler keeps no CUDA link
dependency; on a machine where neither answers, the report prints exactly
what it printed before. The grid itself is a runtime value the compiler
never sees, which is why the report gives the threshold rather than a
launched-work percentage.

The backend half is libmtlc eefb5e6, named in libmtlc.version as always.
Both test suites pass in full (761 of 761, 759 of 759).

Full Changelog: v0.15.0...v0.15.1

Mettle v0.15.0

Choose a tag to compare

@suidvandiewereld suidvandiewereld released this 31 Jul 15:14

The GPU target, hardened by a real workload. An LLM inference engine spent
weeks making Gemma decode fast on Mettle's PTX path and filed eight feature
requests against the language. This release lands all eight.

Warp-per-row kernels can return early

The natural shape for a high-occupancy matvec is a multi-warp block with one
row per warp:

kernel(block = 256) matvec(w: float32*, x: float32*, out: float32*,
                           d: int32, n: int32) {
  var row: int32 = block.x * 8 + thread.x / 32;
  if (row >= d) { return; }          // accepted now
  ...
  var total: float32 = subgroup_reduce_add(sum);
}

The uniformity verifier used to reject that early return. It knew thread.x
varies per lane, so everything derived from it varied too, and kernels had to
clamp the row and carry a live flag instead. The verifier now recognizes
thread.x / 32, thread.x >> 5, and thread.x / subgroup_size() as
subgroup-uniform: every lane of a warp computes the same row, so the guard is
uniform for every warp that reaches the collective. Guards that really do
vary per lane are still rejected. And that rejection is now a plain source
diagnostic; it used to escalate into an internal compiler error that named an
unrelated function.

Kernels declare their launch shape

kernel(block = 256), or kernel(block = (x, y, z)), records the block
geometry a kernel was written for. PTX stamps it as .reqntid and SPIR-V as
LocalSize, so a launch with any other shape fails at the driver with an
error code. Before, a host launching [d, 32] against an 8-warp kernel read
garbage from seven of its eight warps and nothing diagnosed it. Kernels
without the attribute keep their any-geometry behavior.

Bit patterns and packed bytes

Two intrinsic families for custom number formats, declared as externs like
h2f/f2h:

extern fn f32_from_bits(bits: uint32) -> float32 = "f32_from_bits";
extern fn bits_from_f32(x: float32) -> uint32 = "bits_from_f32";
extern fn dp4a_u32(a: uint32, b: uint32, c: uint32) -> uint32 = "dp4a_u32";
extern fn dp4a_s32(a: int32, b: int32, c: int32) -> int32 = "dp4a_s32";
  • f32_from_bits / bits_from_f32 reinterpret a float32 and its IEEE-754
    encoding in either direction: one mov.b32 in PTX, one OpBitcast in
    SPIR-V. Assembling fp8, microscaling, or next year's format no longer means
    arithmetic reconstruction.
  • dp4a_u32 / dp4a_s32 compute the four-way packed-byte dot product with a
    32-bit accumulate, a0*b0 + a1*b1 + a2*b2 + a3*b3 + c. PTX emits the native
    dp4a instruction, collapsing the shift/mask/convert/FMA chain that
    dominates quantized decode; SPIR-V replays the exact byte semantics in
    scalar code.

@unroll(n)

The GPU backends never unroll a loop on their own; the docs now say so
plainly. When a latency-bound inner loop wants its loads pipelined, annotate
it:

@unroll(4) while (j < n) {
  sum = sum + w[row * n + j] * x[j];
  j = j + 32;
}

The compiler emits a main loop that runs four bodies per trip and keeps the
original loop as the remainder, so iteration order, count, and side effects
are preserved exactly for every trip count. It applies to counted loops with
straight-line bodies and a constant positive step, in kernels and in CPU code
under -O. Loops outside that shape stay rolled: the annotation is a hint,
not a contract.

dispatch ... on stream

The compact launch form can now name a stream:

dispatch matvec[(d + 7) / 8, 256](w, x, out, d, n) on stream;

Sugar over the named stream: control, which already existed. It lets a
one-line launch overlap with the previous token's asynchronous readback
without spelling the full three-dimensional form.

--report-occupancy

With --emit-ptx, the compiler runs ptxas -v on the module it just wrote
and prints each kernel's registers per thread plus the occupancy ceiling they
imply, tightened to whole blocks when the kernel declares its shape:

Occupancy report (sm_121a; upper bound: 64K regs/SM, 48 warps/SM, ...):
  matvec: 25 registers, block 256 (8 warps/block, 6 blocks) -> 48/48 resident warps (100%)
  attention: 40 registers -> 48/48 resident warps (100%)

The one-warp-block starvation the engine shipped with for a day would have
been one line in CI.

16-bit typed loads: already there, now pinned

The request asked for uint16* indexing that lowers to ld.global.u16; it
turned out to work already, including the 2-aligned block offsets of GGML's
Q6_K. Regression tests now pin both directions on both backends, and the docs
state the guarantee.

Validation

  • libmtlc suite: 761 of 761. Mettle suite: 759 of 759. Every new construct
    round-trips through ptxas for sm_121a.
  • @unroll was checked bit-exact against the rolled loops across trip counts
    0 through 40, odd strides, < and <=, and every build mode.
  • The requesting engine's own build passes unchanged with this compiler:
    quantization decoders, GPU-versus-CPU logit verification across five model
    architectures with zero argmax mismatches, generation smoke tests, and
    batched prefill equivalence.

The backend half of this work is libmtlc ecb8557, named in
libmtlc.version as always.

Full Changelog: v0.14.2...v0.15.0

v0.14.2

Choose a tag to compare

@github-actions github-actions released this 30 Jul 21:10

Full Changelog: v0.14.1...v0.14.2

v0.14.1

Choose a tag to compare

@github-actions github-actions released this 30 Jul 16:39

Full Changelog: v0.14.0...v0.14.1

Mettle v0.14.0

Choose a tag to compare

@github-actions github-actions released this 30 Jul 13:19

Mettle stops carrying its own backend. It compiles against libmtlc, and this
release names the exact backend commit it uses.

BREAKING CHANGE

Building from source now needs libmtlc.
Fetch it once, then build as before:

./get-libmtlc.sh     # Windows: .\get-libmtlc.ps1
make                 # Windows: .\build.bat

Installing a release needs nothing extra. The backend is already in the binary.

The backend source has left this repository: src/codegen, src/linker,
src/compiler, src/debug, and the IR core with its optimizer. Send patches
for any of those to libmtlc.

Where each half lives

  • This repository holds the language. The lexer, the parser, the type checker,
    the memory safety analysis, IR lowering, the driver, the runtime, the
    standard library.
  • libmtlc holds the IR, the optimizers, code generation for x86-64, ARM64, PTX
    and SPIR-V, and PE and ELF linking.
  • libmtlc.version names one commit, so you can rebuild the compiler that
    shipped. Set LIBMTLC_DIR to build against a checkout of your own.
  • docs/mettle-and-libmtlc.md covers the split, the
    include paths, and how to work on both halves at once.

Fixed: two silent miscompiles

The differential fuzzer found both. --verify missed both, and called the
affected program clean.

  • A shared scaled index addressed the wrong element. a[i] * b[i] computes
    i << 2 twice, and once the optimizer folds the two into one value, two
    address computations read it. The x86 address fold wanted that value to have
    a single reader, so the pair fell through to a fallback that folded the
    already-scaled value in as a unit index. Release builds on the register
    allocating backend read the wrong element. The fold now handles a shared
    index and keeps the shift for whoever else needs it, which is also one
    instruction fewer than before.
  • A spilled base and a spilled index could land in one register. The scaled
    load and address encoders staged each spilled operand through a scratch
    register, and for one combination of allocations they picked the same one, so
    [base + index*4] encoded as [r11 + r11*4]. Each staging register is now
    chosen clear of the ones already in use.

Faster compiles

  • Profiling on 200k-line inputs found the compiler in linear strcmp scans, a
    getenv per pass event, and heap churn. Hashed lookups, lazy teardown and a
    leaner inliner take a 226k-line build from 2244ms to 1082ms, and peak memory
    from 892MB to 508MB.
  • --verify costs 2 to 4 times a plain compile, down from 8.5. The interpreter
    no longer zeroes 570KB per machine, and the pass driver reuses its snapshot
    when a pass changed nothing. ui_demo.mettle drops from 6.5s to 0.36s.

The ML optimizer now proves its work

  • --ml-opt runs every model decision through the reference interpreter before
    it stands. On a divergence it rolls back, then re-applies one decision at a
    time, so the bad one is named with a counterexample and dropped while the
    rest hold. No --verify needed.
  • That makes --ml-opt-speculative usable. The model's dead-code deletions
    carry no proof, so they stand on the validator's word alone.
  • Across the 49 benchmark programs: 1,190 proposals, 348 applied, 537 rejected
    with counterexamples. Every binary matched its baseline at run time.
  • Writing the speculative action found four holes in the validator. This closes
    all four. Each had been blinding --verify too.

--explain-json: stable ids and four new sections

  • Every decision carries a short id beside its prose. Wordings improve;
    int32-sum-narrow-acc does not. Tools should read the ids.
  • A per-function table of instruction counts before and after.
  • A pass ledger that says what each pass did and which lines it touched. The
    vectorizer's row reports both the SIMD kernel it added and the scalar work it
    retired.
  • Per-loop cycles per iteration and the port each loop bottlenecks on. The
    optimizer knows which loops it refused to vectorize; only code generation
    knows what those loops then cost.
  • A call graph and a hotspot ranking.
  • docs/explain-json.md documents the schema.

Mettle for CLion (new)

A JetBrains plugin in tools/clion-plugin. It loads in CLion and across the
IntelliJ family.

  • Diagnostics come from the compiler, so the editor shows the same codes and
    help lines as the command line.
  • The debugger drives --debug-hooks: breakpoints, stepping, the call stack,
    and variables you can edit. Windows only for now.
  • An optimization report window reads --explain-json and applies the fixes the
    compiler already checked.
  • Lexer, parser, completion, navigation and rename cover the language as the
    compiler parses it, down to optional end of line semicolons.

Other

  • The benchmark harness builds its C baselines with clang under -Clang.
  • Every CI job fetches the pinned backend and caches it. The sanitizer job
    instruments the backend too, which is most of the compiler.

What's Changed

Full Changelog: v0.13.0...v0.14.0