nki(linear_self_attention): NKI (Trainium) implementation

[bowencui123]

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

Files: A benchmarks/operators/linear_self_attention/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 (57a84ac)

NKI tunables	Triton counterpart	note
stageA_block_size_m/d, stageB_block_size_m/d (two tuners)	BLOCK_M/BLOCK_N/BLOCK_K/BLOCK_D	—

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
[linear_self_attention] default : verify=OK (2s) 
[linear_self_attention] autotune: verify=OK (0s) last_config={'stageA_block_size_m': 128, 'stageA_block_size_d': 128, 'stageB_block_size_m': 128, 'stageB_block_size_d': 128} trace_records=2 
STUB_EXIT=0
Params    |    Dtype |  Torch(ms) |      NKI(ms) |  Speedup(N)
n=32768   | fp32     |     0.0377 |       0.0616 |        0.61