Skip to content

t27c PTX backend: ternary x int8 GEMM kernel emitted by a .t27 spec (int8 mma.sync), ptx-run glue - #10425

Merged
gHashTag merged 18 commits into
masterfrom
claude/blissful-wescoff-a3c8cd
Oct 19, 2026
Merged

t27c PTX backend: ternary x int8 GEMM kernel emitted by a .t27 spec (int8 mma.sync), ptx-run glue#10425
gHashTag merged 18 commits into
masterfrom
claude/blissful-wescoff-a3c8cd

Conversation

@gHashTag

@gHashTag gHashTag commented Oct 18, 2026 •
edited
Loading

Copy link
Copy Markdown
Owner

Closes #10404. Plan: #10360 (owner override 2026-10-18: build the backend, and build it in t27).

What this adds

  • specs/igla/coder/ternary_gemm_ptx.t27 writes the PTX of one kernel, ternary_gemm_s8: the GPU lowering of specs/igla/coder/ternary_gemm_gpu.t27 on int8 tensor cores.
    • It uses mma.sync.aligned.m16n8k32.row.col.s32.s8.s8.s32 (sm_80+, the RTX 4090 is sm_89). It does not use dp4a.
    • Each block computes a 128x128 tile of C; the k-step is 32 and there are 256 threads, 8 warps of 64x32, so 16 mma per k-step.
    • The 2-bit weight codes are decoded once per tile load, as (c & 1) - ((c >> 1) & 1), and stored transposed as s8 in shared memory.
    • Shared tiles use a 48-byte row stride, so every fragment load is free of bank conflicts.
    • v1 has no cp.async and no double buffering.
  • The same spec holds the host side, so the bit-exact check is t27 too:
    • gen_a and gen_wp, the data;
    • ref_row, the CPU reference;
    • mismatches, the count;
    • shape_ok and grid_dim.
  • 21 tests pin the text: the header, the entry signature, exact mma lines, 16 mma, 64 accumulator inits, 24 fragment loads with their offsets, 16 decodes, 32 guarded stores, the loop shape, and the exact length (21403 bytes). They also check that a short buffer reports the full length, that output is deterministic, the decode table, the packing round trip, the reference on hand cases and the i32 extremes, the mismatch counting (stride and last row), shapes and grid. 5 invariants: i32 overflow bound, shared memory size, stride, warp coverage, 500m shapes.
  • Thin Rust glue. There is no new foreign file, and Rust builds no PTX:
    • t27c compile specs/igla/coder/ternary_gemm_gpu.t27 --backend ptx writes the emitter's bytes to ternary_gemm_gpu.ptx. Every other spec is refused: this lowers one kernel and is not a general t27 -> PTX lowering.
    • t27c ptx-run exists only behind the new cargo feature cuda. It loads the PTX with the CUDA driver API, runs it, copies C back, and calls the spec's mismatches. The default build does not link libcuda.
    • Budget: +70 hand-written foreign lines (main.rs 28, host/driver.rs 39, Cargo.toml 3). The deletion of the orphan bootstrap/src/math_pslq.rs (589 lines, never compiled; docs/reports/orphan_modules.json ceiling 6 -> 5) pays for them.

Verification

  • Railway lab: t27c test-report gives 21/21 pass, 0 vacuous, 5 invariants proved. The seal is saved (gen_hash_rust equals the committed bootstrap/gen/rust/igla/coder/ternary_gemm_ptx.rs).
  • RTX 4090 (RunPod, CUDA 12.4, head c9a20a08c), results comment:
    • The emitted PTX is 21403 bytes, sha256 9fd1ccbc...e96c.
    • ptxas -arch=sm_89: 105 registers, 0 spills.
    • 0 mismatches against the t27 CPU reference on every shape, from 1x16x16 and the ragged 300x272x96 up to 16384x3456x1280.
    • Throughput is 135-174 TOPS. That matches CUTLASS bf16 on the same pod and is 2.5-3.1x below CUTLASS s8. CUTLASS stands in for torch; a Python one-liner is not allowed.
    • The next lever is cp.async double buffering with a k-tile of 64.
  • CI: three checks fail for reasons outside this PR. t27b-native-linux and t27b-native-macos fail on a master @rem test in cli/t27b/tests/tail.rs, which this PR does not touch. claw-sarif fails because the App token permissions are not granted.

🤖 Generated with Claude Code

gHashTag and others added 5 commits October 18, 2026 18:16
@gHashTag @claude
igla ptx: spec that writes the PTX of the ternary x int8 mma.sync GEM…
422d5ce 
…M kernel, with its host reference (Refs #10404)

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@gHashTag @claude
igla ptx: comments inside fn bodies are // (Refs #10404)
cce2fe5 
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@gHashTag @claude
igla ptx: pin the emitted length, fix the store-offset count (Refs #1…
772be75 
…0404)

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@gHashTag @claude
t27c: --backend ptx and ptx-run (cuda feature) as thin glue over the …
36b674b 
…t27 emitter; delete orphan math_pslq.rs to pay the budget (Refs #10404)

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@gHashTag @claude
seal: igla-coder-ternary-gemm-ptx, 21/21 tests, 0 vacuous (Refs #10404)
87ea11f 
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@gHashTag gHashTag added the no-auto-merge Opt out of the issue gate's default squash auto-merge (issue-gate.yml) label Oct 18, 2026
This was referenced Oct 18, 2026
Merged
Merged
@gHashTag @claude
igla ptx: one writer call per statement, so gen-rust output borrows o…
bd2f632 
…ut once; driver: split the truncate borrow (Refs #10404)

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
This was referenced Oct 18, 2026
Merged
Merged
@gHashTag @claude
gen + seal: igla-coder-ternary-gemm-ptx regenerated after the writer …
c9a20a0 
…split, 21/21 (Refs #10404)

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
This was referenced Oct 18, 2026
Merged
Merged
Merged
Merged
Draft
@gHashTag

gHashTag commented Oct 18, 2026

Copy link
Copy Markdown
Owner Author

RTX 4090 run: the t27-emitted PTX is bit-exact; its speed matches bf16 and is about 2.7x below CUTLASS s8

Head c9a20a08c. One RunPod pod: secure cloud, US-NC-1, NVIDIA GeForce RTX 4090, 3105 MHz, 450 W, CUDA 12.4. The pod was deleted after the run.

What ran

  • t27c compile specs/igla/coder/ternary_gemm_gpu.t27 --backend ptx -o ternary_gemm_s8.ptx wrote 21403 bytes, sha256 9fd1ccbca592213c11f8f3544b7e40206245e940698872079d9ba2c06c88e96c. That is the length the spec test pins.
  • ptxas -arch=sm_89 -v: 105 registers, 12288 bytes smem, 0 bytes stack, 0 spill stores, 0 spill loads.
  • t27c ptx-run (--features cuda) loads that PTX with the driver API, launches it, copies C back, and counts mismatches with the spec's own mismatches against ref_row. Both functions are in t27.

Bit-exact

m x n x k rows checked mismatches
1 x 16 x 16 every row 0
128 x 128 x 32 every row 0
300 x 272 x 96 (ragged, every guard path) every row 0
4096 x 1280 x 1280 every row 0
4096 x {3456,1280} x {1280,3456} every 7th 0
16384 x {3456,1280} x {1280,3456}, 16384 x 1280 x 1280 every 31st 0

Throughput, same pod, ms per call

The torch columns are filled by CUTLASS 3.5.1 cutlass_profiler, built on the pod for sm_89. Each column is the fastest kernel the profiler found for that shape. A Python torch one-liner is not allowed here, and CUTLASS is the reference torch itself calls.

m x n x k t27 PTX (s8 IMMA, 2-bit decode) CUTLASS bf16 CUTLASS s8 t27 / s8
4096 x 3456 x 1280 0.2211 (164 T) 0.2237 (162 T) 0.0892 (407 T) 2.5x
4096 x 1280 x 3456 0.2538 (143 T) 0.2219 (163 T) 0.0811 (447 T) 3.1x
4096 x 1280 x 1280 0.0997 (135 T) 0.0901 (149 T) 0.0398 (338 T) 2.5x
16384 x 3456 x 1280 0.8487 (171 T) 0.8490 (171 T) 0.3160 (459 T) 2.7x
16384 x 1280 x 3456 0.8345 (174 T) 0.8430 (172 T) 0.2650 (547 T) 3.1x
16384 x 1280 x 1280 0.3359 (160 T) 0.3183 (169 T) 0.1298 (414 T) 2.6x

How each column was measured, so the two are not mistaken for one method:

  • t27: ptx-run runs one warm-up launch, then times iters=50 back-to-back launches between Instant::now() and cuCtxSynchronize(), so the wall clock includes launch overhead.
  • CUTLASS: CUDA events, 10 warm-up and 50 profiled iterations.
  • The t27 kernel reads 2-bit packed weights and decodes them in the kernel. CUTLASS s8 reads weights that are already s8. That is 4x the weight bytes, but this is not a bandwidth-bound regime at these shapes.

What this says

  • The backend exists and is ours: the kernel text comes from specs/igla/coder/ternary_gemm_ptx.t27, Rust only loads and launches it, and the result is bit-exact.
  • v1 is as fast as CUTLASS bf16 and 2.5-3.1x slower than CUTLASS s8. On the 500m shapes, int8 does not pay off over bf16 until the kernel gets the standard pipeline.
  • The CUTLASS s8 winners are 128x128_64x5 and 128x256_64x3: k-tile 64, 3-5 stages, cp.async. v1 has k-tile 32, one stage, synchronous loads and a bar.sync pair per k-step. The next lever is cp.async with double buffering and k-tile 64. Fragment loads via ldmatrix come after that. Each lands as a spec change with the emitted text pinned again.

Raw outputs are in the session notes (pod_bench.out, pod_prof.out, pod_csv.out).

@gHashTag
gHashTag marked this pull request as ready for review October 18, 2026 11:53
@github-actions github-actions Bot mentioned this pull request Oct 18, 2026
Merged
@github-actions

github-actions Bot commented Oct 18, 2026

Copy link
Copy Markdown
Contributor

PR Dashboard

Generated at: 2026-10-18 12:02:58 UTC

Summary

Status Count
Total Open PRs 50
PRs with Failing Checks 49
PRs with All Checks Green 1
READY 0
FAILING 49
PENDING 0
NO CHECKS YET 0

These columns do not partition: 0 + 49 + 0 + 0 = 49, and there are 50 open PRs. A PR is being counted twice or not at all.

Seal Status

  • ⚠️ STALE -- sha256(compiler.rs)=fad12492937a != manifest seal=87e5cbd3ad94.
    The committed NMSE numbers were certified against an older compiler.rs.
    Run scripts/reseal-check.sh locally for the two-step reseal command (advisory; not a merge gate).

This was referenced Oct 18, 2026
Closed
Merged
Merged
Merged
Merged
@gHashTag gHashTag mentioned this pull request Oct 18, 2026
Draft
This was referenced Oct 18, 2026
Merged
Merged
This was referenced Oct 19, 2026
Merged
Merged
Closed
@gHashTag @claude
t27b ledger: list three specs master added unlisted, all t27b pass as…
f3628bd 
… CI measured (Refs #10404)

specs/fpga/analog_cell_ax7203.t27 and trios-igla-race bin/debug_queue.t27,
bin/smoke_agent.t27.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@github-actions github-actions Bot mentioned this pull request Oct 19, 2026
Merged
@github-actions

github-actions Bot commented Oct 19, 2026

Copy link
Copy Markdown
Contributor

PR Dashboard

Generated at: 2026-10-19 02:41:01 UTC

Summary

Status Count
Total Open PRs 50
PRs with Failing Checks 40
PRs with All Checks Green 10
READY 2
FAILING 40
PENDING 0
NO CHECKS YET 0

These columns do not partition: 2 + 40 + 0 + 0 = 42, and there are 50 open PRs. A PR is being counted twice or not at all.

Seal Status

  • ⚠️ STALE -- sha256(compiler.rs)=fad12492937a != manifest seal=87e5cbd3ad94.
    The committed NMSE numbers were certified against an older compiler.rs.
    Run scripts/reseal-check.sh locally for the two-step reseal command (advisory; not a merge gate).

This was referenced Oct 19, 2026
Merged
Merged
Merged
Open
Merged
Merged
Merged
Merged
Merged
Merged
Merged
Merged
Merged
Merged
Merged
Merged
Merged
Merged
Merged
Merged
@gHashTag
gHashTag merged commit 535334d into master Oct 19, 2026
44 of 46 checks passed
This was referenced Oct 19, 2026
Closed
Merged
Closed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

No reviews

Assignees

No one assigned

Labels

no-auto-merge Opt out of the issue gate's default squash auto-merge (issue-gate.yml)

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

t27c PTX backend: ternary x int8 GEMM kernel emitted by a .t27 spec, run on RTX 4090 (int8 mma.sync)

1 participant

@gHashTag