.. meta::
    :description: Complete reference of environment variables for Neuron Explorer profiling, capture, and configuration.
    :date-modified: 10/19/2026

.. _neuron-explorer-env-vars:

Neuron Explorer Environment Variables
=====================================

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

.. contents:: On this page
   :local:
   :depth: 2


Profiling control
-----------------

These variables enable and configure profiling at runtime.

.. list-table::
   :header-rows: 1
   :widths: 35 45 10 10

   * - 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_NUM_WARMUP``
     - Number of executions to skip before device profiling begins. When set
       to N, the first N executions of each model are skipped and only the
       (N+1)-th execution is profiled. Only supported with model mode
       (``NEURON_RT_INSPECT_DEVICE_PROFILE=1``). Ignored for other modes.
     - ``0``
     - ``0`` – ``4294967295``

**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).

**NEURON_RT_INSPECT_DEVICE_PROFILE_NUM_WARMUP usage:**

Use this variable to bypass warmup noise and profile steady-state behavior.
For example, setting ``NEURON_RT_INSPECT_DEVICE_PROFILE_NUM_WARMUP=5`` skips the
first 5 executions of each model and profiles the 6th. Each model's execution
count is tracked independently — if two different models are loaded, their
warmup counters are separate.

.. note::

   This variable is only effective when ``NEURON_RT_INSPECT_DEVICE_PROFILE=1``
   (model mode). If used with session mode or when device profiling is disabled,
   a warning is logged and the value is ignored.


Source code and debug info
--------------------------

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

.. list-table::
   :header-rows: 1
   :widths: 35 50 15

   * - Variable
     - Description
     - Default
   * - ``NEURON_FRAMEWORK_DEBUG``
     - Also writes the compiled NEFF and the HLO module to the working
       directory. Applies to XLA-based workloads. Does not affect what the
       profile contains.
     - ``0`` (off)
   * - ``NKI_DEBUG_INFO``
     - Enables NKI kernel debug info capture. Produces
       ``kernel_debug_info.json`` in the NEFF. This file provides both the
       instruction-to-NKI-source mapping and the kernel-level NKI Call Stack
       shown in Event Details.
     - ``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 or the NKI Call Stack 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
-------------------------

.. list-table::
   :header-rows: 1
   :widths: 40 45 15

   * - Variable
     - Description
     - Default
   * - ``NEURON_RT_ENABLE_DGE_NOTIFICATIONS``
     - Enables Descriptor Generation Engine notifications so DMA transfers are labeled
       with their tensor variable name (instead of ``unknown``).
     - ``0`` (off)
   * - ``NEURON_RT_PROFILE_BUF_H2D_DMA_MB``
     - Device memory, in MB, reserved for recording host-device DMA events
       in device profiles. Events after the buffer fills are dropped. Set to
       ``0`` to turn off host-device DMA capture. See
       :ref:`Host-device DMA tracks <neuron-explorer-host-device-dma>`.
     - ``64``

.. 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.

.. list-table::
   :header-rows: 1
   :widths: 40 45 15

   * - 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:**

.. code-block:: bash

   # 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:**

.. code-block:: bash

   # 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
-------------------------------

.. list-table::
   :header-rows: 1
   :widths: 40 45 15

   * - 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
----------------

.. list-table::
   :header-rows: 1
   :widths: 40 45 15

   * - 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
------------

.. list-table::
   :header-rows: 1
   :widths: 40 45 15

   * - 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
---------------------

.. list-table::
   :header-rows: 1
   :widths: 40 45 15

   * - 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):**

.. code-block:: bash

   export NEURON_RT_INSPECT_ENABLE=1
   export NEURON_RT_INSPECT_OUTPUT_DIR=./output

**System + device profile (full picture):**

.. code-block:: bash

   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):**

.. code-block:: bash

   # 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):**

.. code-block:: bash

   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):**

.. code-block:: bash

   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