[AntonOresten]

The scalar .b128 ledger assumed LLVM's NVPTX inline asm has no 128-bit register constraint. So each value crossed the asm boundary as two UInt64 words, bridged through block-local .b128 registers with mov.b128 pack/unpack written into the asm text. The constraint does exist: q allocates a PTX .b128 register for an i128 operand on sm_70+, the floor of every .b128 form. It's in the pinned backend (checked on LLVM 22 and on libnvptx 23.1.1).

Changes

• B128 is now an alias for UInt128 (previously NTuple{2,UInt64}), and constraint_letter(UInt128) = "q".
• b128(lo, hi) and b128(x0, x1, x2, x3) pack words low-first into the UInt128, the lane order of mov.b128 d, {lo, hi}. b128(::UInt128) is the identity.
• The B128Ledger emitter names .b128 operands directly and drops the declaration/setup/teardown bridge. For example, ld.global.b128 goes from =l,=l,l with a wrapped asm block to ld.global.b128 $0, [$1]; with =q,l.
• The transpiler emits b128(...) for both mov.b128 pack widths (previously a bare tuple for the two-.b64 form).
• Docs: the dsl.md .b128 section and the mov note in SURFACE.toml.

Emitted PTX

The instruction sequence is unchanged; the pack/unpack moves out of the asm string into LLVM's own code:

main:   { .reg .b128 b128_result; ld.global.b128 b128_result, [%rd1]; mov.b128 {%rd4, %rd5}, b128_result; }
branch: ld.global.b128 %rq2, [%rd1];
        mov.b128 {%rd5, %rd6}, %rq2;

Values do not stay in %rq registers between asm blocks. i128 is still legalized into two i64 in between, so a load feeding a store round-trips through %rd pairs as it did before. This change simplifies the representation; it isn't a codegen win.

Below sm_70, libnvptx reports Inline asm with 128 bit operands is only supported for sm_70 and higher! as a compile diagnostic, not a process abort.

Breaking

Code that built a B128 as (lo, hi) or destructured one as a tuple now uses b128(lo, hi) and integer ops (x % UInt64, (x >> 64) % UInt64).

Verification

Full local suite on the GB10 (CC 12.1, CUDA 13.4): 237,975 passed, 0 failed. That includes gpu/b128_load_store (runtime round-trip of both words) and ptxas/b128_forms (sm_75/80/90/100 partitions).

Independent of #172; the two touch disjoint lines.

🤖 Generated with Claude Code

[codecov]

Codecov Report

✅ All modified and coverable lines are covered by tests.

📢 Thoughts on this report? Let us know!