Mettle v0.15.0

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