nki(matmul_int8): NKI (Trainium) implementation

[bowencui123]

NKI (AWS Trainium) implementation of matmul_int8, split out of the consolidated NKI branch cecilia/feature/nki-vector-add (nki-all-operators, #259) so each operator can be reviewed independently. Supersedes #181 (older per-operator branch: legacy neuronxcc.nki imports; this is the migrated import nki version).

Files: A benchmarks/operators/matmul_int8/impl_nki.py

Status: imports and exposes run()/get_last_config() on trn2 (nki 0.6.0); not individually re-benchmarked in this split

Implementation by @Cecilia123li. Timing/identity infrastructure: #261; Trainium peak/roofline infra: #262.

🤖 Generated with Claude Code

https://claude.ai/code/session_012Q38kGmXvyoeM1qtCbheSL

Autotune (2083bdb)

NKI tunables	Triton counterpart	note
block_size_m, block_size_k, block_size_n	BLOCK_SIZE_M/BLOCK_SIZE_N/BLOCK_SIZE_K	block_size_k counts unpacked K (4 x packed rows)

autotune=False keeps the previous constants (default numbers unchanged). Validation on trn2, case 0 (default run + autotune code path with the candidate timer stubbed — no sweep; --autotune runs a real sweep):

# initial run
[matmul_int8] default : verify=OK (3s) 
[matmul_int8] autotune: verify=OK (9s) last_config={'block_size_m': 128, 'block_size_k': 512, 'block_size_n': 128} trace_records=1 
STUB_EXIT=0
Params         |    Dtype |  Torch(ms) |      NKI(ms) |  Speedup(N)
n=8589934592   | int8     |     0.9921 |       0.2966 |        3.35

[bowencui123]

Review pass (Bowen's assignment): Triton/cuTile alignment + scaffold removal + real autotune & default runs on trn2 (post-#309 runtime-trace timing framework, case 0, --warmup 2 --repeat 10; autotune numbers come from a real selector sweep replayed by the profile worker).

#289 matmul_int8 — NO CODE CHANGE

• Approach ✓ mirrors Triton's packed 2-bit unpack (4 fields, mask+scale→bf16, exact fp32 PSUM
  accumulation → int32). block_size_m/k/n aligned.
• Tested (M2048 N2048 K1024): default 0.2938 ms (torch 0.9879, 3.36×), autotune
  0.2940 ms — 7 candidates timed (0.292–1.669 ms), winner 512/1024/512 (= default);
  non-divisible/non-compiling candidates skipped as designed.

🤖 Generated with Claude Code

https://claude.ai/code/session_011Wvs1tztZTGQFD78YdZaha