This document is relevant for: Inf1, Inf2, Trn1, Trn2, Trn3

Neuron Explorer Environment Variables

This page documents all environment variables that affect Neuron Explorer profiling, profile capture, and profile processing.

Profiling control

These variables enable and configure profiling at runtime.

Variable	Description	Default	Values
NEURON_RT_INSPECT_ENABLE	Master switch to enable runtime profiling. Required for all environment-variable-based capture.	0 (off)	0, 1
NEURON_RT_INSPECT_OUTPUT_DIR	Directory where profile output is written.	./output	Any valid path
NEURON_RT_INSPECT_SYSTEM_PROFILE	Capture system-level runtime events and operations.	1 (on when INSPECT_ENABLE=1)	0, 1
NEURON_RT_INSPECT_DEVICE_PROFILE	Capture device-level hardware traces.	0 (off)	0, 1/model, session

NEURON_RT_INSPECT_DEVICE_PROFILE modes:

• 0 — Disabled (system profile only)

• 1 or model — Captures first execution per NEFF per core (synchronous). Best for compiled-graph workloads.

• session — Captures all device activity in a single NTFF file. Use for async workloads, continuous traces, or inference serving (e.g., vLLM).

Source code and debug info

These variables must be set before compilation to enable source-level mapping in the profiler.

Variable	Description	Default
NEURON_FRAMEWORK_DEBUG	Enables HLO-level stack trace capture. Required for the Source Code Viewer to map hardware instructions to Python model code.	0 (off)
NKI_DEBUG_INFO	Enables NKI kernel source location capture. Produces kernel_debug_info.json in the NEFF for instruction-to-NKI-source mapping.	False
XLA_IR_DEBUG	Enables XLA IR debug info for richer operation names in profiles.	0 (off)
XLA_HLO_DEBUG	Enables XLA HLO debug info for descriptive layer names in the hierarchy view.	0 (off)

Note

NKI_DEBUG_INFO=True is not enabled by default because it increases compile time and memory consumption. Only enable when you need NKI source code linking in the profiler.

Warning

These variables affect compilation output. You must recompile after setting them — existing cached NEFFs won’t have the debug info.

DGE and DMA configuration

Variable	Description	Default
NEURON_RT_ENABLE_DGE_NOTIFICATIONS	Enables Data Gather Engine notifications so DMA transfers are labeled with their tensor variable name (instead of unknown).	0 (off)

Warning

Enabling DGE notifications adds overhead and can cause timeouts for very large NEFFs. If profiling hangs or times out, disable this flag and re-capture. The profile will be complete but DMA variable names will show unknown.

Capture filtering (capture-time)

These reduce memory and file size by filtering events at capture time. Data not captured is permanently lost.

Variable	Description	Default
NEURON_RT_INSPECT_EVENT_FILTER_NC	Only capture events from specified NeuronCores. Accepts comma-separated indices or ranges.	All cores
NEURON_RT_INSPECT_EVENT_FILTER_TYPE	Only capture specified event types or categories.	All events
NEURON_RT_INSPECT_SYS_TRACE_MAX_EVENTS_PER_NC	Maximum system trace events per NeuronCore before ring buffer overwrites oldest. Increase if you see dropped events.	1,000,000

NeuronCore filter examples:

# Only NeuronCore 0
export NEURON_RT_INSPECT_EVENT_FILTER_NC=0

# NeuronCores 0, 2, and 4
export NEURON_RT_INSPECT_EVENT_FILTER_NC=0,2,4

# Range: NeuronCores 0 through 3
export NEURON_RT_INSPECT_EVENT_FILTER_NC=0-3

Event type filter examples:

# Only hardware events
export NEURON_RT_INSPECT_EVENT_FILTER_TYPE=hardware

# Only software events
export NEURON_RT_INSPECT_EVENT_FILTER_TYPE=software

# Hardware events except cc_exec
export NEURON_RT_INSPECT_EVENT_FILTER_TYPE=hardware,^cc_exec

# Specific event types
export NEURON_RT_INSPECT_EVENT_FILTER_TYPE=model_load,nrt_execute,runtime_execute

Event type groups:

• hardware: nc_exec_running, cc_running, cc_exec_barrier, numerical_err, nrt_model_switch, timestamp_sync_point, hw_notify

• software: All other events

NEFF cache and artifact control

Variable	Description	Default
TORCH_NEURONX_NEFF_CACHE_DIR	Directory for the NEFF compilation cache. When set, cached NEFFs are automatically available in the profile output directory.	System default
NEURON_RT_INSPECT_PRECACHE_ENABLE	Enable profiling before graph cache is populated. Useful for capturing the first execution of a model.	0 (off)

PyTorch-specific

Variable	Description	Default
TORCH_NEURONX_ENABLE_HOST_CC	Enable host-mediated collective communication for distributed profiling.	0 (off)
TORCH_NEURONX_ENABLE_ASYNC_NRT	Enable asynchronous Neuron Runtime for overlapping compute and communication.	0 (off)

Warning

Do not set NEURON_RT_INSPECT_ENABLE=1 and use the PyTorch profiling APIs (for example, torch.profiler) simultaneously. They conflict. Use one or the other.

JAX-specific

Variable	Description	Default
NEURON_RT_INSPECT_DEVICE_PROFILE	For JAX, set to 1 with jax.profiler.trace for device captures. Do not combine with NEURON_RT_INSPECT_ENABLE=1.	0 (off)

Warning

Do not set both NEURON_RT_INSPECT_ENABLE=1 and use jax.profiler.trace simultaneously. They conflict. Use one or the other.

Runtime configuration

Variable	Description	Default
NEURON_RT_NUM_CORES	Number of NeuronCores to use for the workload.	All available
NEURON_RT_VIRTUAL_CORE_SIZE	Number of physical NeuronCores per virtual NeuronCore (VNC).	1

Quick reference: recommended combinations

System profile only (lightweight):

export NEURON_RT_INSPECT_ENABLE=1
export NEURON_RT_INSPECT_OUTPUT_DIR=./output

System + device profile (full picture):

export NEURON_RT_INSPECT_ENABLE=1
export NEURON_RT_INSPECT_DEVICE_PROFILE=1
export NEURON_RT_INSPECT_OUTPUT_DIR=./output

Full debug info (system + device + source code + NKI + DMA names):

# Set BEFORE compilation
export NEURON_FRAMEWORK_DEBUG=1
export NKI_DEBUG_INFO=True
export XLA_IR_DEBUG=1
export XLA_HLO_DEBUG=1

# Set BEFORE execution
export NEURON_RT_INSPECT_ENABLE=1
export NEURON_RT_INSPECT_DEVICE_PROFILE=1
export NEURON_RT_ENABLE_DGE_NOTIFICATIONS=1
export NEURON_RT_INSPECT_OUTPUT_DIR=./output

Inference / vLLM (session-based device profile):

export NEURON_RT_INSPECT_ENABLE=1
export NEURON_RT_INSPECT_DEVICE_PROFILE=session
export NEURON_RT_INSPECT_OUTPUT_DIR=./output

Large distributed workload (filtered to reduce size):

export NEURON_RT_INSPECT_ENABLE=1
export NEURON_RT_INSPECT_DEVICE_PROFILE=1
export NEURON_RT_INSPECT_EVENT_FILTER_NC=0,1
export NEURON_RT_INSPECT_OUTPUT_DIR=./output

This document is relevant for: Inf1, Inf2, Trn1, Trn2, Trn3