nki.language.device_print#

nki.language.device_print(print_prefix, tensor)[source]#

Print a message with a string print_prefix followed by the value of a tile tensor.

By default, using this function will not result in your tensors being printed out. When running your kernel, you need to define the environment variable NEURON_RT_DEBUG_OUTPUT_DIR and point it to a directory that will store the tensor data grouped by prefix each time the device_print instruction is executed.

The structure of the directory will be <print_prefix>/core_<logical core id>/<iteration>/....

Listing 12 Example usage#
import nki.isa as nisa
import nki.language as nl

def my_nki_kernel(input_tensor):
    a_tile = sbuf.view(input_tensor.dtype, input_tensor.shape)
    nisa.dma_copy(a_tile, input_tensor)
    nl.device_print("a_tile", a_tile)

    ...

Warning

This feature is only available when using the NxD Inference library.

Parameters:

  • print_prefix (str) – prefix of the print message. This string is evaluated at trace time and must be a constant expression.

  • tensor – tensor to print out. Can be in SBUF or HBM.

Returns:

None