Skip to content

nki(quantize_global): NKI (Trainium) implementation - #295

Open
bowencui123 wants to merge 6 commits into
mainfrom
bowen/nki/quantize_global
Open

bowencui123 wants to merge 6 commits into
mainfrom
bowen/nki/quantize_global

Conversation

@bowencui123

@bowencui123 bowencui123 commented Aug 29, 2026

Copy link
Copy Markdown
Collaborator

NKI (AWS Trainium) implementation of quantize_global, split out of the consolidated NKI branch cecilia/feature/nki-vector-add (nki-all-operators, #259) so each operator can be reviewed independently.

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

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
[quantize_global] default : verify=OK (1s) 
[quantize_global] autotune: verify=OK (0s) last_config={'block_size': 16384} trace_records=1 
STUB_EXIT=0
Params      |    Dtype |  Torch(ms) |      NKI(ms) |  Speedup(N)
n=1048576   | fp32     |     0.0264 |       0.0619 |        0.43

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.

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
bowencui123 and others added 5 commits August 29, 2026 08:37
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
…NC2)

Merges NKI backend timing into results/csv/quantize_global_default.csv, run
against this branch's impl_nki.py on trn2.3xlarge with the LNC2
execution contract (NEURON_LOGICAL_NC_CONFIG=2, NEURON_RT_NUM_CORES=1,
NEURON_CC_FLAGS="--target trn2 --lnc 2").

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01AQseF7nyesBh8KZAp8g7Cm
…e, LNC2)

Merges NKI backend timing into results/csv/quantize_global_autotune.csv, run with
--autotune against this branch's impl_nki.py on trn2.3xlarge, LNC2
execution contract. All cases pass correctness verification.

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01AQseF7nyesBh8KZAp8g7Cm
…obal

# Conflicts:
#	results/csv/quantize_global_autotune.csv
#	results/csv/quantize_global_default.csv
…pad/reshape); rerun default+autotune benchmarks

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01ScXYNjrrKGgDUVNHxv7HJt
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