This document is relevant for: Trn1, Trn2, Trn3

nki.language.no_reorder#

nki.language.no_reorder()[source]#

Prevent the scheduler from reordering operations in this region.

Use as a context manager (with nl.no_reorder():) to guarantee that operations inside the block execute in program order. Without this directive, the compiler scheduler is free to reorder independent operations for better hardware utilization.

Dynamic loops (nl.dynamic_range) are not supported inside a no_reorder block. Static loops (nl.affine_range, nl.sequential_range, nl.static_range) are allowed because they are fully unrolled at compile time.

Examples:

import nki.language as nl

# nki.language.no_reorder -- guarantee execution order
with nl.no_reorder():
    a = nl.full((128, 512), 3.0, dtype=nl.float32, buffer=nl.sbuf)
    b = nl.full((128, 512), 2.0, dtype=nl.float32, buffer=nl.sbuf)
    c = nl.add(a, b)
expected = nl.full((128, 512), 5.0, dtype=nl.float32, buffer=nl.sbuf)
assert nl.equal(c, expected)

This document is relevant for: Trn1, Trn2, Trn3