Skip to content

Conversation

@jlonge4
Copy link
Contributor

@jlonge4 jlonge4 commented Nov 5, 2025

Issue #, if available:

N/A

Description of changes:

Demonstrates batch invariance behavior in NKI by replicating findings from Thinking Machines' "Defeating Nondeterminism in LLM Inference".

Key findings:

  1. Replicated the paper: nki.lang kernels (MatMul, RMSNorm) show batch variance in both float32 and bfloat16 when using dynamic reduction strategies
  2. New discovery: nki.isa operations show variance in float32 but no variance in bfloat16, making them more suitable for deterministic inference

Contents:

  • MatMul and RMSNorm kernel implementations (both lang and ISA variants)
  • Test suite comparing batch variance across kernel types and precisions
  • Documentation explaining findings and practical implications

Testing:

Please see detailed unit test requirements in the CONTRIBUTING.md

  • The change is covered by numeric check using nki.baremetal
  • The change is covered by performance benchmark test using nki.benchmark
  • The change is covered by end-to-end integration test

Pull Request Checklist

  • I have filled in all the required field in the template
  • I have tested locally that all the tests pass
  • By submitting this pull request, I confirm that my contribution is made under the terms of the MIT-0 license.

jlonge4 and others added 18 commits October 10, 2025 15:04
Refactor tests for batch invariance and variance in RMSNorm and MatMul. Now follows the same testing pattern as Thinking Machines Labs.
Added tests for matmul kernel correctness and batch variance effects. Updated existing tests to improve clarity and structure.
Updated RMSNorm kernel to demonstrate batch variance with split reduction along the hidden dimension. Adjusted tile sizes based on batch invariance parameter to illustrate the impact on floating-point addition order during reduction.
Refactor RMSNorm kernel tto replace nl.arange with nl.mgrid
Replaced direct load/store operations with nisa.dma_copy for better performance.
Using DMA copy for improved performance.
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