nki(1d_conv): NKI (Trainium) implementation

[bowencui123]

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

Files: A benchmarks/operators/1d_conv/impl_nki.py

Status: imports and exposes run()/get_last_config() on trn2 (nki 0.6.0); not individually re-benchmarked in this split; imports via importlib (the operator name starts with a digit, so a plain import statement is invalid — the engine uses importlib)

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

🤖 Generated with Claude Code

https://claude.ai/code/session_012Q38kGmXvyoeM1qtCbheSL

Autotune (98ffb58)

NKI tunables	Triton counterpart	note
block_size_l (output-length block), blocks_in_flight	BLOCK_SIZE_BATCH_LENGTH/BLOCK_SIZE_IN_FEAT/BLOCK_SIZE_OUT_FEAT	feature blocks are the 128-partition tiles

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
[1d_conv] default : verify=FAIL (2s) Comparing
[1d_conv] autotune: verify=FAIL (1s) last_config={'block_size_l': 2048, 'blocks_in_flight': 2} trace_records=1 Comparing
STUB_EXIT=1
Params          |    Dtype |  Torch(ms) |      NKI(ms) |  Speedup(N)
n=12884901888   | fp16     |     0.1437 |       0.4109 |        0.35
# re-validation (tolerance parsing fixed in the checker; engine line above stands)
[1d_conv] default : verify=OK (1s) 
[1d_conv] autotune: verify=OK (0s) last_config={'block_size_l': 2048, 'blocks_in_flight': 2} trace_records=1 
STUB_EXIT=0

[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).

#263 1d_conv — NO CODE CHANGE

• im2col-as-affine-views + matmul ≙ Triton's im2col approach; block_size_l/blocks_in_flight
  tunables; SBUF-budget-driven defaults; dynamic-loop interior. Clean, no scaffold found.
• Tested (b1 C128→128 k3 L=131072 fp16): default 0.4167 ms (torch 0.1595), autotune
  0.3862 ms — 6 candidates timed (0.342–0.675 ms baremetal), winner block_size_l=4096, blocks_in_flight=4 beats the SBUF-heuristic default (8192, 4).

🤖 Generated with Claude Code

https://claude.ai/code/session_011Wvs1tztZTGQFD78YdZaha