nki(dequantize_rowwise): NKI (Trainium) implementation

[bowencui123]

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

Files: A benchmarks/operators/dequantize_rowwise/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 (ed94192)

NKI tunables	Triton counterpart	note
block_size (free-dim tile per 128-partition block)	BLOCK_SIZE	—

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
[dequantize_rowwise] default : verify=FAIL (1s) Comparing
[dequantize_rowwise] autotune: verify=FAIL (0s) last_config={'block_size': 8192} trace_records=1 Comparing
STUB_EXIT=1
Params     |    Dtype |  Torch(ms) |      NKI(ms) |  Speedup(N)
n=262144   | int8     |     0.0250 |       0.0247 |        1.01
# re-validation (tolerance parsing fixed in the checker; engine line above stands)
[dequantize_rowwise] default : verify=OK (0s) 
[dequantize_rowwise] autotune: verify=OK (0s) last_config={'block_size': 8192} trace_records=1 
STUB_EXIT=0