This document is relevant for: Trn2, Trn3

nki.isa.NkiInstruction

class nki.isa.NkiInstruction(handle)

Opaque handle to a traced NKI instruction.

Returned by nki.isa operations during tracing. Use .depends_on(*producers) to add scheduling dependencies (this instruction runs after every listed producer completes).

Methods

depends_on

Add scheduling dependencies: self (consumer) runs after each producer.

depends_on(*producers: NkiInstruction) → NkiInstruction

Add scheduling dependencies: self (consumer) runs after each producer.

Returns self, so the two forms compose:

c.depends_on(b, a)             # one call, two edges
c.depends_on(b).depends_on(a)  # chained, same effect

A dependency is not a trace-time hint: it lowers to a hardware semaphore wait that the device executes every time this instruction runs. Because it is fixed at compile time and fires unconditionally, both instructions must sit at the same dynamic-loop nesting (or the producer in an enclosing scope). You cannot wire an edge into or out of a dynamic loop (nl.fori_loop / nl.while_loop), or its semaphore could fire unconditionally or never fire (deadlock). Static Python for/if are unrolled at trace time and do not count.

This document is relevant for: Trn2, Trn3