Repository navigation
igla-coder: ternary x int8 GPU GEMM kernel spec, bit-exact in three modes (Closes #10361) - #10362
Merged
Merged
Conversation
…odes (Closes #10361) specs/igla/coder/ternary_gemm_gpu.t27 lays the kernel out as CUDA would: block_tile is the __global__ body for one block, with zero-padded shared tiles and per-thread register micro-tiles; weights are packed 2-bit and unpacked at tile load; gemm_tiled is the grid launch. Three inner-product modes (add/subtract, i8 MMA, dp4a) are compared with == against a scalar reference. That is valid because products lie in [-128,128] with i32 accumulation, so nothing overflows while k <= 16777215 (an invariant). The spec imports nothing: the race references it mirrors do not compile under t27c or t27b today; the correct 8x8 path is copied in. Railway lab, 2026-10-18: t27c 17/17 tests, 7 invariants proved, 0 vacuous; t27b (aarch64 JIT) 17/17, 7 held, also under --overflow trap. A negative control with two mutated assertions fails exactly those two under both compilers. Seal from t27c seal --save on the lab; dupe scan reports no new duplicate body. Refs #10360 (cost model and gen-cuda label request). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
gHashTag
enabled auto-merge (squash)
October 18, 2026 09:31
Contributor
This was referenced Oct 18, 2026
Closed
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Closes #10361. Refs #10360, the RFC with the cost model and the gen-cuda
owner-approved-foreignlabel request.What
specs/igla/coder/ternary_gemm_gpu.t27is a tiled ternary x int8 GEMM laid out the way a CUDA kernel is:block_tileis the__global__body for one block, with zero-padded shared tiles and per-thread register micro-tiles.gemm_tiledis the grid launch..trinity/seals/coder_igla-coder-ternary-gemm-gpu.jsonis the seal, written byt27c seal --saveon the Railway lab.Every tiling and every mode is compared with
==against a scalar reference. Bit-exact comparison is valid because products lie in [-128, 128] and accumulation is i32. An invariant proves nothing overflows while k <= 16777215; d_ff = 3456 is far below that.The spec imports nothing.
specs/igla/race/ternary_gemm.t27andternary_mac.t27do not compile under t27c or t27b today, and they contain wrong argument orders and a wrong expected value. Details are in #10360. The correct 8x8 path is mirrored here instead.Evidence (Railway lab, 2026-10-18, spec sha256 573a8240...6d)
--overflow trap: also 17 of 17.No foreign code was added. The backend and the bench wait for the label on #10360.
🤖 Generated with Claude Code