Skip to content

nki(rope): NKI (Trainium) implementation - #300

Open
bowencui123 wants to merge 2 commits into
mainfrom
bowen/nki/rope
Open

bowencui123 wants to merge 2 commits into
mainfrom
bowen/nki/rope

Conversation

@bowencui123

@bowencui123 bowencui123 commented Aug 29, 2026

Copy link
Copy Markdown
Collaborator

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

Files: A benchmarks/operators/rope/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 (079265f)

NKI tunables Triton counterpart note
block_size (sequence rows per tile, <=128) 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
[rope] default : verify=FAIL (4s) Comparing
[rope] autotune: verify=FAIL (0s) last_config={'block_size': 128} trace_records=1 Comparing
STUB_EXIT=1
Params      |    Dtype |  Torch(ms) |      NKI(ms) |  Speedup(N)
n=4194304   | fp16     |     0.2257 |       0.9116 |        0.25
# re-validation (tolerance parsing fixed in the checker; engine line above stands)
[rope] default : verify=OK (0s) 
[rope] autotune: verify=OK (0s) last_config={'block_size': 128} trace_records=1 
STUB_EXIT=0

Split out of the consolidated NKI branch cecilia/feature/nki-vector-add
(nki-all-operators, PR #259) so each operator can be reviewed on its own.
Supersedes PR #189 (older per-operator branch).

Co-Authored-By: Cecilia123li <68335867+Cecilia123li@users.noreply.github.com>
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_012Q38kGmXvyoeM1qtCbheSL
Tunables mirror the Triton search space (`BLOCK_SIZE`); defaults are the previous constants,
so autotune=False is unchanged.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_012Q38kGmXvyoeM1qtCbheSL
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant