Skip to content

nki(softmax): NKI (Trainium) implementation - #302

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

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

Conversation

@bowencui123

@bowencui123 bowencui123 commented Aug 29, 2026

Copy link
Copy Markdown
Collaborator

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

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

Status: verified on trn2 (default + autotune; NKI 0.055 ms @ 2048x512 fp16)

  • includes NkiAutotuner wiring (search space over the kernel's tile/loop parameter; defaults unchanged)

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

🤖 Generated with Claude Code

https://claude.ai/code/session_012Q38kGmXvyoeM1qtCbheSL

Autotune (0ac2225)

NKI tunables Triton counterpart note
free_cap (one-pass vs streaming switch), block_size (streamed column tile) BLOCK_SIZE/BLOCK_N

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
[softmax] default : verify=OK (1s) 
[softmax] autotune: verify=OK (0s) last_config={'free_cap': 8192, 'block_size': 2048} trace_records=1 
STUB_EXIT=0
Params      |    Dtype |  Torch(ms) |      NKI(ms) |  Speedup(N)
n=1048576   | fp16     |     0.0404 |       0.0551 |        0.73

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 #191 (older per-operator branch).
- includes `NkiAutotuner` wiring (search space over the kernel's tile/loop parameter; defaults unchanged)

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`/`BLOCK_N`); 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.

2 participants