.. NOTE: This page is generated by explorer/scripts/generate-schema-docs.py in the neuron-explorer source repository. Do not edit it by hand, edit the schema YAML or that script and regenerate. .. meta:: :description: Reference for the parquet table schemas produced by Neuron Explorer, matching ``neuron-explorer view --show-profile-schema``. .. _neuron-explorer-profile-schema-reference: Profile Parquet Schema Reference ================================ When Neuron Explorer processes a profile it writes the results as a set of `Parquet `_ tables (``Summary.parquet``, ``Instruction.parquet``, ``DmaPacket.parquet``, and so on). This page documents every table and field in that output to assist in interpreting all the profile data. The profile schema is defined as an OpenAPI document that is embedded in the ``neuron-explorer`` binary. The :ref:`Schema Reference ` section below is generated directly from that definition, so it matches exactly what ``neuron-explorer view --show-profile-schema`` prints. This page reflects schema version ``12.0.1`` released in version **2.33.0** of the Neuron SDK. Getting the schema for your installed version --------------------------------------------- The schema evolves with each Neuron release. To print the exact schema for the version you have installed: .. code-block:: bash neuron-explorer view --show-profile-schema This prints the OpenAPI schema as YAML, one block per table, covering both device-profile and system-profile tables. The schema is built into the `neuron-explorer` binary and can be accessed without any profile. If you have a profile, you can find out which schema version a profile was produced with by reading the ``profile_schema_version`` field of the :ref:`Metadata ` table. The version uses ``MAJOR.MINOR.PATCH`` semantics: ``MAJOR`` is an incompatible change (a table or field removed/renamed, or a type change), ``MINOR`` is a compatible addition (a new table or field), and ``PATCH`` is a metadata-only change (such as an improved description). Profile output format ---------------------- Parquet is the default and recommended output format. It is columnar and compactly encoded, so producing and querying it is fast even for large profiles, and the Neuron Explorer UI reads from the same Parquet tables, so what you query matches what the UI shows. The legacy ``--output-format json`` output is less complete, less organized, and less scalable, and so is not recommended for new analysis. How the profile data is organized --------------------------------- Every table holds one *kind* of data. Knowing a table's kind tells you how to read its rows: .. list-table:: :header-rows: 1 :widths: 22 38 40 * - Kind of data - What one row represents - Example tables * - Timeline of events - Each is a discrete event, with a start and end timestamp (or a single timestamp). - ``Instruction``, ``DmaPacket``, ``DmaPacketAggregated``, ``SemaphoreUpdate``, ``Throttle``, ``Error``, ``CcOp``, ``ActiveTime`` * - Time-series samples - Each is a value at a point in time. - ``DmaUsage``, ``HbmUsage``, ``PsumUsage``, ``SbufUsage``, ``CpuUsage``, ``HostMemUsage`` * - Dependency-graph edges - Each is a directed edge between rows in timeline tables (e.g. an instruction and the DMA it triggered). - ``Flow`` * - Hierarchical aggregation - Each is a node in the compiler IR hierarchy (Framework → HLO → tensorizer → backend → Instruction), with rolled-up statistics. - ``FrameworkInstruction``, ``HloInstruction``, ``PenguinInstruction``, ``BirInstruction``, ``FrameworkNode`` * - Aggregated summary - Each is a summary metric computed across the whole profile (or by some key). - ``Summary``, ``OpcodeSummary``, ``ThrottleSummary``, ``HbmUsageSummaryByType`` * - Reference / lookup - Each is an entity in the model or execution that other tables join to them by id to enrich rows without duplicating the detail. - ``TensorInfo``, ``DmaQueuesInfo``, ``CcStream``, ``StackFrame``, ``KernelInstructions`` * - Profile metadata - Each describes a piece of profile metadata. - ``Metadata``, ``NeffHeader``, ``SystemProfileMetadata``, ``ExecutionInfo`` * - Diagnostics - Each is a warning emitted during ingestion. Worth checking first, rows here can indicate a data-quality problem. - ``Warning`` A few types (the ``MemoryBandwidth*`` types) are *transient*: they are computed at API query time and returned by Explorer's HTTP API after profile processing rather than written to Parquet during profile processing. The table description will note if the table is transient. The output Parquet is also self-describing. After a profile is processed, the ``SchemaFields`` Parquet table contains one row per ``(table, field)`` with its type, description, unit, and other metadata. This ``SchemaFields`` table contains the same information shown on this page, in a queryable format. Units and value ranges ---------------------- The **Unit** column gives each field's unit, and the description ends with the field's declared range where it has one. Fields whose names end in ``_percent`` are, with a few exceptions, **fractions between 0 and 1**, not percentages: ``0.45`` means 45%, and ``1.0`` means 100%. These show **fraction (0-1)** in the Unit column. The Neuron Explorer UI multiplies them by 100 for display, but the stored and printed values are the fractions. The exceptions are labelled explicitly: **percent (0-100)** is already a percentage, and a plain **fraction** has no upper bound and can exceed 1, as its description explains. In the ``SchemaFields`` table, both kinds carry ``field_unit`` ``percent``; read ``field_maximum`` to tell them apart (``1`` for a fraction, ``100`` for a percentage). The ``Summary`` table holds the same metrics, under the same field names and on the same scale, that ``neuron-explorer view --output-format summary-json`` and ``--output-format summary-text`` print, so this page also documents that output. .. _profile-schema-reference-tables: Schema Reference ---------------- The tables and enums below are generated from the Neuron Explorer's OpenAPI schema definition, which describes all profile data. For each field, if the **Type** column references an enum, it will be linked. The **Unit** column gives the measurement unit where one applies. And the **Required** indicates whether the field is always present. **Tables** * :ref:`ActiveTime ` * :ref:`AssemblyInstruction ` * :ref:`BirInstruction ` * :ref:`CcInstruction ` * :ref:`CcOp ` * :ref:`CcStream ` * :ref:`CompilerEnvVars ` * :ref:`CompilerFlags ` * :ref:`CoreBarriers ` * :ref:`CpuUsage ` * :ref:`DeviceProfileList ` * :ref:`DgePhaseGroup ` * :ref:`DmaPacket ` * :ref:`DmaPacketAggregated ` * :ref:`DmaQueuesInfo ` * :ref:`DmaUsage ` * :ref:`DmaUsageSummary ` * :ref:`DtypeSummary ` * :ref:`Error ` * :ref:`ExecutionInfo ` * :ref:`Flow ` * :ref:`FrameworkInstruction ` * :ref:`FrameworkNode ` * :ref:`HbmAllocation ` * :ref:`HbmUsage ` * :ref:`HbmUsageSummaryByType ` * :ref:`HloInstruction ` * :ref:`HostDeviceDmaUsage ` * :ref:`HostMemUsage ` * :ref:`Instruction ` * :ref:`KernelInstructions ` * :ref:`KernelIterationVariables ` * :ref:`KernelStackFrames ` * :ref:`MemoryBandwidthPoint ` * :ref:`MemoryBandwidthResponse ` * :ref:`MemoryBandwidthSeries ` * :ref:`Metadata ` * :ref:`NeffHeader ` * :ref:`NkiKernelCallSites ` * :ref:`OpcodeSummary ` * :ref:`PendingDma ` * :ref:`PenguinInstruction ` * :ref:`PsumUsage ` * :ref:`SbufAllocation ` * :ref:`SbufUsage ` * :ref:`SchemaFields ` * :ref:`SemaphoreUpdate ` * :ref:`SnapshotAllocation ` * :ref:`SnapshotSegment ` * :ref:`StackFrame ` * :ref:`StackFrameFileLocation ` * :ref:`StackFrameFileName ` * :ref:`StackFrameFunctionName ` * :ref:`Summary ` * :ref:`SystemProfileEvents ` * :ref:`SystemProfileHbmUsage ` * :ref:`SystemProfileMetadata ` * :ref:`TensorInfo ` * :ref:`Throttle ` * :ref:`ThrottleSummary ` * :ref:`TimelineAnnotation ` * :ref:`Warning ` **Enums** * :ref:`CaptureType ` * :ref:`DataType ` * :ref:`DependencyType ` * :ref:`DmaQueueType ` * :ref:`ErrorType ` * :ref:`ExecutionUnit ` * :ref:`MemoryBandwidthDirection ` * :ref:`PerformanceMode ` * :ref:`SpillType ` Tables ------ .. _profile-schema-activetime: ActiveTime ~~~~~~~~~~ A timeline of when engines are active. Engines are considered active when it is processing at least one instruction (excluding semaphore waits). This is created by merging the intervals of all instructions on the engine. This is different than the instruction table, which gives individual instructions without merging and includes all details about individual instructions. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``start_ts`` - integer (int64) - nanoseconds - Yes - Timestamp when activity on this engine starts. * - ``end_ts`` - integer (int64) - nanoseconds - Yes - Timestamp when activity on this engine ends. * - ``engine`` - string - - Yes - Compute engine that this active time period corresponds to. Example: ``tensor``. * - ``pcore_idx`` - integer (int64) - - Yes - Index of the physical NeuronCore within a Logical NeuronCore (LNC). A Logical NeuronCore is a group of physical NeuronCores. For LNC2 this field can have a value of either 0 or 1. * - ``duration_ns`` - integer (int64) - nanoseconds - Yes - The duration of the active time period. .. _profile-schema-assemblyinstruction: AssemblyInstruction ~~~~~~~~~~~~~~~~~~~ Assembly instructions from the NEFF. These are lower level instructions than backend-level instructions. Each backend-level instruction can map to multiple assembly instructions. This table is only used by perf_sim profiles, which are performance simulations of the NeuronCore at the schedulers level. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``instruction_id`` - integer (int64) - - Yes - Foreign key reference to the parent Instruction.id. Links this assembly instruction to its parent backend-level instruction. * - ``compiler_opcode`` - string - - Yes - Opcode of the instruction at compile time. The "compiler_opcode" and the runtime "opcode" can be different because certain missing or dynamic information is filled in at runtime. For example, the "compiler_opcode: PSEUDO_DMA_TRIGGER" may change to "opcode: WRITE" at runtime because that is when the write address is known. Example: ``PSEUDO_DMA_TRIGGER``. * - ``compiler_operands`` - string - - Yes - Operands of the instruction at compile time. The "compiler_operands" and the runtime "operands" can be different because certain missing or dynamic information is filled in at runtime. For example, a compiler_opcode PSEUDO_DMA_TRIGGER with "compiler_operands: qSyncPIOParam0 block_id=1" may change to an opcode WRITE with "operands: data=0x1,0x1 size=4,4 dst=0xffff1009038,0xffff1029038" at runtime because that is when the write address is known. Example: ``name=sg0000 args_table_var_id=0``. * - ``compiler_pc`` - integer (int64) - - Yes - The compiler program counter is this instruction index on the given engine. The "compiler_pc" and the runtime "pc" can be different because the compiler can have instructions that are unrolled at runtime. Or the compiler can have a collective communication operation that is fulfilled by multiple runtime instructions. .. _profile-schema-birinstruction: BirInstruction ~~~~~~~~~~~~~~ Backend-level instruction information for displaying hierarchy graph. Parent level is tensorizer level, child level is instruction level. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``id`` - integer (int64) - - Yes - Unique identifier for corresponding assembly level instruction. * - ``raw_bir_id`` - integer (int64) - - Yes - Unique identifier for backend-level instruction. * - ``raw_penguin_id`` - integer (int64) - - Yes - Unique identifier for tensorizer-level instruction. * - ``raw_hlo_id`` - integer (int64) - - Yes - Unique identifier for hlo level instruction. * - ``raw_framework_id`` - integer (int64) - - Yes - Unique identifier for framework level instruction. * - ``bir_name`` - string - - Yes - Backend-level instruction name. Example: ``I-1347-2_sg0000``. * - ``modular_flow_invocation_idx`` - integer (int64) - - Yes - Integer for identifying which call is the current instruction. 0 represent first time calling the function, 1 represent second time calling the function. Example: ``1``. * - ``pcore_idx`` - integer (int64) - - No - Index of the physical NeuronCore within a Logical NeuronCore (LNC). A Logical NeuronCore is a group of physical NeuronCores. For LNC2 this field can have a value of either 0 or 1. * - ``start_ts`` - integer (int64) - nanoseconds - Yes - Minimum start time among all of the corresponding instruction timestamps. Unit is nanoseconds. * - ``end_ts`` - integer (int64) - nanoseconds - Yes - Maximum end time among all of the corresponding instruction timestamps. Unit is nanoseconds. * - ``time_intervals`` - string - - Yes - A list of timestamp intervals after merging all corresponding instruction timestamps from the lower hierarchy level. Each interval is represented as [start, end]. Example: ``[[1, 10], [20, 50]]``. * - ``element_count_per_interval`` - array of integer (int64) - - Yes - Array where each item corresponds to the total number of elements processed by a group of instructions from the lower hierarchy level. Each corresponds to the time interval in the same position in the time_intervals field. Example: ``[1024, 2048, 512]``. * - ``transpose_element_count_per_interval`` - array of integer (int64) - - Yes - Array where each item corresponds to the total number of transpose elements processed by a group of instructions from the lower hierarchy level. Each corresponds to the time interval in the same position in the time_intervals field. Example: ``[256, 512, 128]``. * - ``hbm_transfer_size_per_interval`` - array of integer (int64) - bytes - Yes - Array where each item corresponds to the total number of HBM bytes transferred by a group of instructions from the lower hierarchy level. Each corresponds to the time interval in the same position in the time_intervals field. Example: ``[16384, 32768, 8192]``. * - ``hbm_read_bytes`` - integer (int64) - bytes - Yes - Total number of bytes read from HBM by all grouped instructions in this backend-level instruction. * - ``hbm_write_bytes`` - integer (int64) - bytes - Yes - Total number of bytes written to HBM by all grouped instructions in this backend-level instruction. * - ``sbuf_read_bytes`` - integer (int64) - bytes - No - Total number of bytes read from SBUF by all grouped instructions in this backend-level instruction. * - ``sbuf_write_bytes`` - integer (int64) - bytes - No - Total number of bytes written to SBUF by all grouped instructions in this backend-level instruction. * - ``load_weight_bytes`` - integer (int64) - bytes - No - Total number of bytes loaded by LDWEIGHTS instructions in this backend-level instruction. .. _profile-schema-ccinstruction: CcInstruction ~~~~~~~~~~~~~ Collective Communication (CC) instructions that are executed by the CC-Cores. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``execution_index`` - integer (int32) - - No - Zero-based index of the NEFF execution this instruction belongs to. Matches execution_index in ExecutionInfo, FrameworkNode, NkiKernelCallSites, and Instruction. Use to scope instructions to a specific execution within a multi-execution profile. Null on profiles ingested before this field existed. * - ``start_ts`` - integer (int64) - nanoseconds - Yes - Timestamp when the collective instruction started execution. * - ``end_ts`` - integer (int64) - nanoseconds - Yes - Timestamp when the collective instruction finished execution. * - ``duration_ns`` - integer (int64) - nanoseconds - Yes - Duration of the collective instruction. * - ``core_idx`` - integer (int64) - - Yes - The index of the CC-Core that executed the collective instruction. * - ``description`` - string - - Yes - The name of the collective instruction as provided by NeuronCollectivesFirmware. Example: ``DMA_ADVANCE``. * - ``trigger_engine`` - string - - No - The engine that triggered the collective instruction. Example: ``Pool``. * - ``trigger_start_ts`` - integer (int64) - nanoseconds - No - Start timestamp of the PSEUDO_TRIGGER_COLLECTIVE instruction which triggered this collective instruction. .. _profile-schema-ccop: CcOp ~~~~ Collective Commmunication (CC) operations that are executed by the collective engines. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``id`` - integer (int64) - - Yes - Unique identifier for this collective operation within the profile. Used as in_id or out_id in the Flow table to link the operation to the instruction that triggered it. Derived from stream_id and start_ts, which together identify one row. * - ``start_ts`` - integer (int64) - nanoseconds - Yes - Timestamp when the collective operation started execution. * - ``end_ts`` - integer (int64) - nanoseconds - Yes - Timestamp when the collective operation finished execution. * - ``duration_ns`` - integer (int64) - nanoseconds - Yes - Duration of the collective operation. * - ``operation`` - string - - Yes - The name of the collective operation being performed. Example: ``AllGather``. * - ``algorithm`` - string - - Yes - The algorithm used for performing the collective operation. Example: ``Ring``. * - ``num_elements`` - integer (int64) - - Yes - Number of elements in the collectives operation. * - ``dtype`` - string - - Yes - Data type used for the collective operation. Example: ``fp16``. * - ``op_id`` - integer (int64) - - No - ID of the collective operation. Absent on a barrier row, which is a synchronization marker rather than a collective operation and so has no id. * - ``stream_id`` - integer (int64) - - Yes - Unique ID of the stream that the collective operation is on for a given profile. * - ``stream_index`` - integer (int64) - - Yes - ID of the stream that the collective operation is on relative to the execution. * - ``trigger_engine`` - string - - Yes - The engine that triggered the collective operation. Example: ``Pool``. * - ``input_size`` - integer (int64) - bytes - No - The input size in bytes of the collective operation. * - ``output_size`` - integer (int64) - bytes - No - The output size in bytes of the collective operation. * - ``algorithm_bandwidth`` - number (float) - bytes_per_second - No - The algorithm bandwidth of the collective operation. Calculated as (max(input_size, output_size) / duration)). * - ``bus_bandwidth`` - number (float) - bytes_per_second - No - The bus bandwidth of the collective operation. This number is obtained by applying a formula to the algorithm bandwidth to reflect the true speed of the collective communication operation. The formula is dependant on the operation and number of cores. * - ``trigger_start_ts`` - integer (int64) - nanoseconds - No - Start timestamp of the PSEUDO_TRIGGER_COLLECTIVE instruction which triggered this collective operation. * - ``cc_trigger_start_delay`` - integer (int64) - nanoseconds - No - The time delay between when the collective operation was triggered, and when execution actually started. * - ``replica_group_id`` - integer (int64) - - No - The group ID that identifies which set of participating ranks/devices will communicate during this collective operation. Each group ID corresponds to a specific replica group configuration. * - ``replica_group`` - string - - No - String representation of the participating ranks/devices in this collective operation, formatted as nested arrays showing the communication topology (e.g., "[[0, 1, 2, 3], [4, 5, 6, 7]]" for two groups of 4 devices each). Example: ``[[0, 1, 2, 3]]``. * - ``src_target_pairs`` - string - - No - For explicit permute collectives, the directed source-to-target rank routing as a list of [source_rank, target_rank] pairs (e.g., "[[0, 1], [1, 0]]" means rank 0 sends to rank 1 and rank 1 sends to rank 0). Permute collectives set this instead of replica_group because they route between specific ranks rather than communicating over a symmetric group. Example: ``[[0, 1], [1, 0]]``. .. _profile-schema-ccstream: CcStream ~~~~~~~~ Information about Collective Communication (CC) streams. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``stream_id`` - integer (int64) - - Yes - The id of this collective communication stream. * - ``stream_index`` - integer (int64) - - Yes - ID of the stream that the collective operation is on relative to the execution. * - ``cc_core_ids`` - string - - Yes - The CC-Core ids assigned to this stream. Example: ``[8,9,10,11]``. .. _profile-schema-compilerenvvars: CompilerEnvVars ~~~~~~~~~~~~~~~ Environment variables captured during NEFF compilation. Each row represents a single environment variable with its name, value, and whether it was explicitly set by the user. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``name`` - string - - Yes - The name of the environment variable. Example: ``NEURON_CC_FLAGS``. * - ``value`` - string - - Yes - The current value of the environment variable. Example: ``--retry_failed_compilation``. * - ``set`` - boolean - - Yes - Whether the environment variable was explicitly set by the user (true) or is using a system/default value (false). Example: ``true``. .. _profile-schema-compilerflags: CompilerFlags ~~~~~~~~~~~~~ Compiler flags used during NEFF compilation. Each row represents a single compiler flag with its name, value, default, and whether it was explicitly set by the user. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``name`` - string - - Yes - The name of the compiler flag. Example: ``optimize_level``. * - ``value`` - string - - Yes - The current value of the compiler flag. Example: ``2``. * - ``default`` - string - - Yes - The default value of the compiler flag. Example: ``1``. * - ``set`` - boolean - - Yes - Whether the flag was explicitly set by the user (true) or is using its default value (false). Example: ``true``. .. _profile-schema-corebarriers: CoreBarriers ~~~~~~~~~~~~ Core barrier operations will synchronize a single compute engine across all physical NeuronCores within a single Logical NeuronCore (LNC). These operations ensure data has landed where it needs to be before execution can continue. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``start_ts`` - integer (int64) - nanoseconds - Yes - Timestamp the core barrier operation started. * - ``end_ts`` - integer (int64) - nanoseconds - Yes - Timestamp the core barrier operation completed. * - ``duration_ns`` - integer (int64) - nanoseconds - Yes - Duration of the core barrier operation. * - ``engine_name`` - string - - Yes - The compute engine that this core barrier operation corresponds to. Example: ``GpSimd``. * - ``engine_instance_idx`` - integer (int64) - - Yes - The compute engine index that this core barrier operation corresponds to. Default: ``0``. * - ``id`` - integer (int64) - - Yes - Unique identifier for this core barrier operation. * - ``semaphore`` - integer (int64) - - Yes - Semaphore used to implement this core barrier. * - ``compiler_pc`` - integer (int64) - - Yes - Compiler program counter for this core barrier instruction. * - ``compiler_opcode`` - string - - Yes - Core barrier opcode at compile time. Example: ``PSEUDO_CORE_BARRIER``. * - ``compiler_operands`` - string - - Yes - Operands at compile time specifying semaphore conditions, completion events, and barrier ID. Example: ``$S[45]>=16 $S[8]++@complete semaphore=9 id=1``. * - ``pcore_idx`` - integer (int64) - - No - Index of the physical NeuronCore within a Logical NeuronCore (LNC). A Logical NeuronCore is a group of physical NeuronCores. For LNC2 this field can have a value of either 0 or 1. .. _profile-schema-cpuusage: CpuUsage ~~~~~~~~ Host CPU utilization time-series metrics. Each row represents a CPU utilization sample for either a specific CPU core or the average across all cores on an instance. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``timestamp`` - integer (int64) - nanoseconds - Yes - Timestamp for this CPU utilization measurement. * - ``util_percent`` - number (float) - percent (0-100) - Yes - CPU utilization percentage at this timestamp. Each CPU utilization value represents the average CPU utilization over the previous 1 second time period. Range: 0 to 100. * - ``cpu_id`` - string - - Yes - Identifier for the CPU core this measurement applies to. Can be a numeric CPU ID (e.g., "0", "1") or "average" for the average utilization across all cores. Example: ``["0", "1", "average"]``. * - ``instance_id`` - string - - Yes - Identifier for the instance on which this CPU utilization was measured. Example: ``i-0b1ea78ca2865fd32``. * - ``ultraserver_id`` - integer (int64) - - No - Unique UltraServer ID if this measurement was captured on an UltraServer instance. * - ``ultraserver_instance_index`` - integer (int64) - - No - The instance index within the UltraServer on which this measurement was captured. .. _profile-schema-deviceprofilelist: DeviceProfileList ~~~~~~~~~~~~~~~~~ This table contains device profile information including device profile names, instance IDs, process IDs, NeuronCore IDs, session IDs, and model IDs. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``device_profile_name`` - string - - Yes - Name of the device profile. Example: ``neuron_profile_001``. * - ``instance_id`` - string - - Yes - Unique identifier of the server instance that produced this profile. Example: ``i-0b1ea78ca2865fd32``. * - ``process_id`` - integer (int64) - - Yes - Process ID that produced this profile. * - ``nc_id`` - integer (int64) - - Yes - NeuronCore ID associated with this profile. * - ``session_id`` - integer (int64) - - No - Unique session identifier for this profile. Present for session-based profiling. * - ``model_id`` - integer (int64) - - No - Unique model identifier for this profile. Present for model-based profiling. .. _profile-schema-dgephasegroup: DgePhaseGroup ~~~~~~~~~~~~~ A wrapping span that groups the GpSimd phase sub-bars (DGE_ALLGATHER, DGE_DESCS_BUILD) of a single indirect-DGE operation into one box. It spans from the first GpSimd phase entry to the parent indirect-DMA instruction's end, so the related GpSimd timeline spans render inside one box. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``start_ts`` - integer (int64) - nanoseconds - Yes - Timestamp the indirect-DGE phase group started (first GpSimd phase entry). * - ``end_ts`` - integer (int64) - nanoseconds - Yes - Timestamp the indirect-DGE phase group completed (parent indirect-DMA instruction end). * - ``duration_ns`` - integer (int64) - nanoseconds - Yes - Duration of the indirect-DGE phase group. * - ``engine_name`` - string - - Yes - The compute engine the phase group renders on. Always GpSimd for indirect-DGE. Example: ``GpSimd``. * - ``id`` - integer (int64) - - Yes - Unique identifier for this phase group. * - ``opcode`` - string - - Yes - Opcode of the parent indirect-DMA instruction this phase group wraps. Example: ``DMA_INDIRECT``. * - ``pcore_idx`` - integer (int64) - - Yes - Index of the physical NeuronCore within a Logical NeuronCore (LNC). A Logical NeuronCore is a group of physical NeuronCores. For LNC2 this field can have a value of either 0 or 1. .. _profile-schema-dmapacket: DmaPacket ~~~~~~~~~ A DMA packet represents a NeuronCore DMA transfer or a host-device DMA transfer, called H2D DMA by the runtime and driver. A NeuronCore DMA packet typically transfers a portion of a tensor. Its source and destination can be HBM or SBUF. The packet traverses from the source location, over the AXI bus, to a DMA engine, over the AXI bus, and to the destination location. A host-device DMA transfers data between host memory and device HBM accessible to a NeuronCore, in either direction. Common uses include transferring application data between host memory and HBM, fetching instructions, and handling work independently of NeuronCore DMA activity. Host-device DMA uses the same packet representation as NeuronCore DMA. When a matching DmaPacketAggregated row exists, packet_aggregation_id associates the packet with the aggregated transfer. The DmaPacketAggregated row records metadata shared by the packet set, including the source, destination, and variable name. Some DMA packets have no additional metadata in the DmaPacketAggregated table. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``start_ts`` - integer (int64) - nanoseconds - Yes - Timestamp when instruction started execution. * - ``end_ts`` - integer (int64) - nanoseconds - Yes - Timestamp when instruction finished execution. * - ``duration_ns`` - integer (int64) - nanoseconds - Yes - Duration of instruction execution. * - ``pcore_idx`` - integer (int64) - - No - Index of the physical NeuronCore within a Logical NeuronCore (LNC). A Logical NeuronCore is a group of physical NeuronCores. For LNC2 this field can have a value of either 0 or 1. * - ``id`` - integer (int64) - - Yes - Unique identifier for this DMA packet. * - ``packet_idx`` - integer (int32) - - Yes - Index of the DMA packet within the set of DMA packets which fulfill a transfer of one or more tensors. For host-device DMA, this is the packet's ordinal within its queue. * - ``packet_aggregation_id`` - integer (int64) - - No - Unique identifier of the corresponding row in the DmaPacketAggregated table. Each aggregated DMA represents the transfer of one or more tensors fulfilled on a Neuron Device by a set of DMA packets. The corresponding DmaPacketAggregated row contains more information about the packet, including the source, destination, variable name, and other values shared by the aggregated packet set. See the DmaPacketAggregated table for more details. This field is null when the packet has no aggregated DMA, which is the case for host-device DMA packets, packets on a DMA queue the NEFF does not describe, and packets the profiler could not attribute to a specific DGE instruction. * - ``is_host_device_dma`` - boolean - - Yes - Whether this packet uses host-device DMA, called H2D DMA by the runtime and driver. Host-device DMA can transfer application data between host memory and NeuronCore-accessible HBM in either direction, fetch instructions, or handle work independently of NeuronCore DMA activity. Example: ``false``. * - ``engine_idx`` - integer (int64) - - No - Index of the DMA engine that executed this DMA packet. * - ``read_bytes`` - integer (int64) - bytes - Yes - Size in bytes of data read during this operation. * - ``write_bytes`` - integer (int64) - bytes - Yes - Size in bytes of data written during this operation. * - ``transfer_bytes`` - integer (int64) - bytes - Yes - Size in bytes of data transferred during this operation. * - ``throughput`` - number (float) - bytes_per_second - Yes - The throughput of this DMA operation. This is the calculated as the amount of data transferred divided by the transfer duration. * - ``queue_type`` - string (:ref:`DmaQueueType `) - - No - The type of the NeuronCore DMA queue that this packet belongs to. Used to identify the origin, and contents of the DMA packet. ``software_dynamic`` packets are generated by the Gpsimd cores. ``hardware_dynamic`` packets are generated by dedicated DGE hardware. ``instruction``, ``input``, ``output``, and ``data`` packets are generated by the compiler. Host-device DMA packets omit this field because they do not use a NeuronCore DMA queue. Allowed values: ``instruction``, ``input``, ``output``, ``data``, ``software_dynamic``, ``hardware_dynamic``, ``collectives``, ``unknown``. * - ``variable`` - string - - No - Variable or tensor name. For a collectives DMA, this is the collective label '.', or the queue name when no collective matches. Example: ``input_tensor``. * - ``spill_type`` - string (:ref:`SpillType `) - - No - Set when this DMA packet is part of a compiler spill or reload, to 'save' (State Buffer -> HBM temp) or 'reload' (HBM temp -> State Buffer). Inherited from the packet's DMA, derived from source and dest. Allowed values: ``save``, ``reload``. * - ``queue_name`` - string - - No - Queue name where this DMA operation was queued. Example: ``qScalarSpillReload0-Q1-E64``. * - ``queue_idx`` - integer (int64) - - No - Queue index where this DMA operation was queued. * - ``crc`` - integer (int64) - - Yes - Cyclic Redundancy Check value for data integrity verification. Example: ``2847562039``. * - ``packet_id`` - integer (int64) - - No - Hardware-assigned sequence number of this packet within its DMA queue, taken from the DMA engine's start notification. Present only when the hardware reports one. * - ``seq_id`` - integer (int64) - - No - Identifier of the sequencer that issued this DMA, taken from the DMA engine's start notification. Present only when the hardware reports one. .. _profile-schema-dmapacketaggregated: DmaPacketAggregated ~~~~~~~~~~~~~~~~~~~ Each row in this table represents an aggregation of a set DMA packets which together fulfil a transfer of one or more tensors. The duration of an aggregated DMA spans from the first start to the last end of the DMA packets in the set. The byte quantities of an aggregated DMA are the sum of the DMA packets in the set. An aggregated DMA has a unique id and a block_id that is only unique per combination of dma_engine, dma_queue, modular_flow_function, and pcore_idx. The id field in the DmaPacketAggregated table corrosponds to the packet_aggregation_id field in the DmaPacket table, linking each DMA packet to its corresponding aggregated DMA. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``start_ts`` - integer (int64) - nanoseconds - Yes - Timestamp when instruction started execution. * - ``end_ts`` - integer (int64) - nanoseconds - Yes - Timestamp when instruction finished execution. * - ``duration_ns`` - integer (int64) - nanoseconds - Yes - Duration of instruction execution. * - ``id`` - integer (int64) - - Yes - Unique identifier for this aggeregated set of DMA packets. * - ``pcore_idx`` - integer (int64) - - No - Index of the physical NeuronCore within a Logical NeuronCore (LNC). A Logical NeuronCore is a group of physical NeuronCores. For LNC2 this field can have a value of either 0 or 1. * - ``block_id`` - integer (int32) - - Yes - A block_id is used by the compiler to logically identify the transfer of one or more tensors. The transfer is fulfilled on a Neuron Device by a set of DMA packets. A block_id only unique per combination of dma_engine, dma_queue, function, and subgraph. * - ``queue_idx`` - integer (int64) - - Yes - Queue index where this DMA operation was queued. * - ``queue_name`` - string - - Yes - Queue name where this DMA operation was queued. Example: ``qScalarSpillReload0-Q1-E64``. * - ``queue_type`` - string (:ref:`DmaQueueType `) - - No - The type of the DMA queue that this packet belongs to. Used to identify the origin, and contents of the DMA packet. ``software_dynamic`` packets are generated by the Gpsimd cores. ``hardware_dynamic`` packets are generated by dedicated DGE hardware. ``instruction``, ``input``, ``output``, and ``data`` packets are generated by the compiler. Allowed values: ``instruction``, ``input``, ``output``, ``data``, ``software_dynamic``, ``hardware_dynamic``, ``collectives``, ``unknown``. * - ``op`` - string - - No - DMA operation type. Example: ``read``. * - ``variable`` - string - - Yes - Variable or tensor name associated with the data being transferred. Example: ``input_tensor``. * - ``read_bytes`` - integer (int64) - bytes - Yes - Size in bytes of data read during this operation. * - ``write_bytes`` - integer (int64) - bytes - Yes - Size in bytes of data written during this operation. * - ``transfer_bytes`` - integer (int64) - bytes - Yes - Size in bytes of data transferred during this operation. * - ``throughput`` - number (float) - bytes_per_second - Yes - The throughput of this DMA operation in bytes per second. Calculated as the amount of data transferred divided by the transfer duration. * - ``source`` - string - - Yes - Source memory location(s) for the DMA transfer. Example: ``[SB]``. * - ``dest`` - string - - Yes - Destination memory location(s) for the DMA transfer. Example: ``[OUTPUT]``. * - ``spill_type`` - string (:ref:`SpillType `) - - No - Set when this DMA is a compiler spill or reload, to 'save' (State Buffer -> HBM temp) or 'reload' (HBM temp -> State Buffer). Derived from source and dest. Allowed values: ``save``, ``reload``. * - ``semaphore_id`` - string - - Yes - Identifier of semaphore used for synchronization. Example: ``S[10] (qGpSimdPIO0)``. * - ``modular_flow_function`` - string - - No - Modular flow function name and instance associated with this DMA operation. Example: ``sg0000:0``. * - ``read_sbuf_partitions`` - string - - No - State Buffer partitions read from during this DMA operation. Example: ``0,1``. * - ``read_num_sbuf_partitions`` - integer (int32) - - Yes - Number of State Buffer partitions read from during this DMA operation. * - ``read_shape`` - string - - Yes - Multi-dimensional shape read from source memory. Example: ``[1024, 512]``. * - ``read_offset`` - string - - Yes - Byte offset within the source memory region. Example: ``[0x2000]``. * - ``read_steps`` - string - - Yes - Stride of reads from source location. Example: ``[1, 32768]``. * - ``write_sbuf_partitions`` - string - - No - State Buffer partitions written to during this DMA operation. Example: ``0-15``. * - ``write_num_sbuf_partitions`` - integer (int32) - - Yes - Number of State Buffer partitions written to during this DMA operation. * - ``write_shape`` - string - - Yes - Multi-dimensional shape written to destination memory. Example: ``[1024, 512]``. * - ``write_offset`` - string - - Yes - Byte offset within the destination memory region. Example: ``[0x1000]``. * - ``write_steps`` - string - - Yes - The destination stride. Example: ``[1, 32]``. * - ``is_transpose_mode`` - boolean - - No - Indicates if transpose mode is enabled. Example: ``true``. * - ``transpose_element_size`` - string - - No - Element size for transpose operations. Example: ``[2]``. * - ``transpose_shape`` - string - - No - Shape descriptor for transpose operation. Example: ``[1024, 4, 512]``. * - ``transpose_from_shape`` - string - - No - Source shape for transpose operations. Example: ``[4, 1024, 512]``. * - ``transpose_to_shape`` - string - - No - Target shape for transpose operation. Example: ``[1024, 512, 4]``. * - ``transpose_unroll_factor`` - string - - No - Unroll factor for transpose operation. Example: ``[128, 32, 1, 2]``. * - ``trigger`` - integer (int64) - - No - Start timestamp of the instruction that triggered this DMA. This DMA had a dependency on that instruction. * - ``completion`` - integer (int64) - - No - End timestamp of the instruction that triggered this DMA. This DMA had a dependency on that instruction. .. _profile-schema-dmaqueuesinfo: DmaQueuesInfo ~~~~~~~~~~~~~ Information about DMA queues. Each DMA engine has 16 DMA queues. DMA queues exist for each compute engine (ex. Tensor engine) and data type (ex. weights). The items in the queues are DMA descriptors. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``name`` - string - - Yes - The name of the DMA queue. Example: ``qSyncDynamicHW``. * - ``type`` - string - - Yes - The type of DMA queue. Example: ``weights``. * - ``pcore_idx`` - integer (int64) - - Yes - Index of the physical NeuronCore within a Logical NeuronCore (LNC). A Logical NeuronCore is a group of physical NeuronCores. For LNC2 this field can have a value of either 0 or 1. * - ``idx`` - integer (int64) - - Yes - The index of this DMA queue within the DMA engine. * - ``engine_idx`` - integer (int64) - - Yes - The index of the DMA engine that this DMA queue belongs to. * - ``axi_port`` - integer (int64) - - Yes - The AXI port used by this DMA queue. * - ``dram_channel`` - integer (int64) - - Yes - The DRAM channel used by this DMA queue. .. _profile-schema-dmausage: DmaUsage ~~~~~~~~ DMA time-series usage metrics such as throughput over time. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``timestamp`` - integer (int64) - nanoseconds - Yes - Timestamp for this measurement. * - ``current_throughput`` - number (float) - bytes_per_second - Yes - Rate of data transfer over DMA in bytes per second at this timestamp. * - ``total_transfer_bytes`` - integer (int64) - bytes - Yes - Total bytes transferred over DMA including transfers to/from HBM and State Buffer. This total is cumulative from the start of profiling to this timestamp. * - ``pcore_idx`` - integer (int64) - - No - Index of the physical NeuronCore within a Logical NeuronCore (LNC). A Logical NeuronCore is a group of physical NeuronCores. For LNC2 this field can have a value of either 0 or 1. .. _profile-schema-dmausagesummary: DmaUsageSummary ~~~~~~~~~~~~~~~ Pre-bucketed DMA usage metrics for fast query-time aggregation. Written during ingestion at a fixed bucket resolution. Internal optimization table — not queried directly by users. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``timestamp`` - integer (int64) - nanoseconds - Yes - Bucket start timestamp aligned to bucket boundary. * - ``pcore_idx`` - integer (int64) - - Yes - Index of the physical NeuronCore. * - ``sum_current_throughput`` - number (double) - - Yes - Sum of current_throughput values in this bucket. Divide by count for average. * - ``sum_total_transfer_bytes`` - integer (int64) - - Yes - Sum of total_transfer_bytes values in this bucket. Divide by count for average. * - ``max_total_transfer_bytes`` - integer (int64) - - Yes - Maximum total_transfer_bytes value in this bucket. * - ``count`` - integer (int64) - - Yes - Number of raw rows aggregated into this bucket. Minimum: 1. * - ``bucket_resolution_ns`` - integer (int64) - nanoseconds - Yes - The bucket resolution in nanoseconds used during ingestion. Minimum: 1. .. _profile-schema-dtypesummary: DtypeSummary ~~~~~~~~~~~~ Summary of instruction counts and execution time grouped by compute engine and the data types of an instruction's operands. Each row represents one (engine, src1_dtype, src2_dtype, dst_dtype) combination. Source data types are stored positionally (src1 = first input operand, src2 = second input operand); single-input instructions leave src2_dtype unset. Consumers derive higher-level views from these rows, such as the source/destination data type pair and the marginal distributions per source and per destination data type. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``engine`` - string - - Yes - Compute engine that executed the instructions. Example: ``tensor``. * - ``src1_dtype`` - string (:ref:`DataType `) - - No - Data type of the first input operand for instructions in this group. Allowed values: ``fp64``, ``fp32``, ``fp32r``, ``tf32``, ``fp16``, ``bf16``, ``fp8e3``, ``fp8e4``, ``fp8e5``, ``mxfp8``, ``mxfp4``, ``int8``, ``int16``, ``int32``, ``int64``, ``uint8``, ``uint16``, ``uint32``, ``uint64``, ``unknown``. Example: ``bf16``. * - ``src2_dtype`` - string (:ref:`DataType `) - - No - Data type of the second input operand for instructions in this group. Unset for single-input instructions. Allowed values: ``fp64``, ``fp32``, ``fp32r``, ``tf32``, ``fp16``, ``bf16``, ``fp8e3``, ``fp8e4``, ``fp8e5``, ``mxfp8``, ``mxfp4``, ``int8``, ``int16``, ``int32``, ``int64``, ``uint8``, ``uint16``, ``uint32``, ``uint64``, ``unknown``. Example: ``bf16``. * - ``dst_dtype`` - string (:ref:`DataType `) - - No - Data type of the output operand for instructions in this group. Allowed values: ``fp64``, ``fp32``, ``fp32r``, ``tf32``, ``fp16``, ``bf16``, ``fp8e3``, ``fp8e4``, ``fp8e5``, ``mxfp8``, ``mxfp4``, ``int8``, ``int16``, ``int32``, ``int64``, ``uint8``, ``uint16``, ``uint32``, ``uint64``, ``unknown``. Example: ``fp32``. * - ``instruction_count`` - integer (int64) - - Yes - Total number of instructions with this engine and data type combination. * - ``instruction_time_ns`` - integer (int64) - nanoseconds - Yes - Total cumulative execution time of all instructions with this engine and data type combination. Overlapping instructions on the same engine are added together (not merged). * - ``active_time_ns`` - integer (int64) - nanoseconds - Yes - Duration of time when at least one instruction with this engine and data type combination is executing. Overlapping instructions are merged so that concurrent execution is counted only once. * - ``active_time_percent`` - number (float) - - Yes - Percentage of total execution time when at least one instruction with this engine and data type combination is executing. Range: 0 to 1. .. _profile-schema-error: Error ~~~~~ Information about hardware generated error events such as Out Of Bounds (OOB) or Not A Number (NAN). .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``description`` - string - - Yes - Short description of the error condition. Example: ``["EMBD OOB", "PDMAC OOB", "DMA2D OOB", "DMAI OOB", "Debug", "Tensor Completion"]``. * - ``error_type_id`` - integer (int32) - - Yes - Numerical value of error type. Distinguishes between different error conditions. Example: ``1``. * - ``metadata`` - string - - No - Error metadata as string. Example: ``sem 5 overflow``. * - ``metadata_hex`` - string - - No - Hexadecimal representation of error metadata. Example: ``1292``. * - ``pcore_idx`` - integer (int64) - - No - Index of the physical NeuronCore within a Logical NeuronCore (LNC). A Logical NeuronCore is a group of physical NeuronCores. For LNC2 this field can have a value of either 0 or 1. * - ``timestamp`` - integer (int64) - nanoseconds - Yes - Timestamp when the error occurred. * - ``type`` - string (:ref:`ErrorType `) - - Yes - Category or type of error. Allowed values: ``Custom Notification``, ``Errors``, ``Warnings``. .. _profile-schema-executioninfo: ExecutionInfo ~~~~~~~~~~~~~ Per-execution information for profiles with one or more executions. Each row corresponds to a single execution within the profile. In multi-NEFF profiles, different executions may correspond to different NEFFs. Use execution_start_ts and execution_end_ts to correlate the time range of each execution's data (instructions, DMAs, etc.) with the NEFF file name. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``execution_index`` - integer (int32) - - Yes - Zero-based index of this execution within the profile. * - ``neff_name`` - string - - No - File name of the NEFF executed during this execution. Nil when the execution has no associated NEFF (e.g. out-of-NEFF collective communication). Example: ``model.neff``. * - ``execution_start_ts`` - integer (int64) - nanoseconds - No - Start timestamp of this execution relative to the start of the profile. Nil for pre-v6 NTFF profiles that lack an execution timeline. * - ``execution_end_ts`` - integer (int64) - nanoseconds - No - End timestamp of this execution relative to the start of the profile. Nil for pre-v6 NTFF profiles that lack an execution timeline. * - ``num_subgraphs`` - integer (int32) - - No - Number of subgraphs processed in this execution. A subgraph is a directory in the unpackaged NEFF that contains files describing the computation and resources needed to run on a single NeuronCore. Each subgraph declared in the NEFF is loaded onto its own NeuronCore. On post-INF1 architectures, multiple subgraphs in a NEFF correspond to the logical core feature. * - ``row_count`` - integer (int64) - - No - Number of rows written to the profile database for this execution. .. _profile-schema-flow: Flow ~~~~ Storing dependency information for rows in different or same table (such as between instructions and DMAs). .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``in_id`` - integer (int64) - - Yes - Unique identifier for the input flow. * - ``out_id`` - integer (int64) - - Yes - Unique identifier for the output flow. * - ``in_table`` - string - - Yes - Table name for the input flow. Example: ``instruction``. * - ``out_table`` - string - - Yes - Table name for the output flow. Example: ``dma_block``. * - ``dependency_type`` - string (:ref:`DependencyType `) - - No - Kind of dependency edge from the input row to the output row. Absent when the kind is unknown. Allowed values: ``ORDERED``, ``ANTI_DEPENDENCE``, ``OUTPUT_DEPENDENCE``, ``FLOW_DEPENDENCE``, ``trigger``, ``completion``, ``DGE_TRIGGER``, ``SAME_PSUM_BANK``, ``INVALID``. * - ``gap_time`` - integer (int64) - nanoseconds - No - The gap time between the end of the input event and the start of the output event. .. _profile-schema-frameworkinstruction: FrameworkInstruction ~~~~~~~~~~~~~~~~~~~~ Framework level instruction information. Instruction in this table has same format as its corresponding debuginfo. Child level is hlo level. For example, layer[0]/layer[1] as one entry of framework level debuginfo would also be created as one FrameworkInstruction. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``id`` - integer (int64) - - Yes - Unique identifier for corresponding assembly level instruction. * - ``raw_bir_id`` - integer (int64) - - Yes - Unique identifier for backend-level instruction. * - ``raw_penguin_id`` - integer (int64) - - Yes - Unique identifier for tensorizer-level instruction. * - ``raw_hlo_id`` - integer (int64) - - Yes - Unique identifier for hlo level instruction. * - ``raw_framework_id`` - integer (int64) - - Yes - Unique identifier for framework level instruction. * - ``modular_flow_invocation_idx`` - integer (int64) - - No - Integer for identifying which call is the current instruction. 0 represent first time calling the function, 1 represent second time calling the function. Example: ``1``. * - ``framework_name`` - string - - No - Framework operation name. Example: ``NeuronLlamaDecoderLayer[1]/function[2]/aten.add``. * - ``start_ts`` - integer (int64) - nanoseconds - Yes - Minimum start time among all of the corresponding instruction timestamps. Unit is nanoseconds. * - ``end_ts`` - integer (int64) - nanoseconds - Yes - Maximum end time among all of the corresponding instruction timestamps. Unit is nanoseconds. * - ``time_intervals`` - string - - No - A list of timestamp intervals after merging all corresponding instruction timestamps from the lower hierarchy level. Each interval is represented as [start, end]. Example: ``[[1, 10], [20, 50]]``. * - ``element_count_per_interval`` - array of integer (int64) - - Yes - Array where each item corresponds to the total number of elements processed by a group of instructions from the lower hierarchy level. Each corresponds to the time interval in the same position in the time_intervals field. Example: ``[1024, 2048, 512]``. * - ``transpose_element_count_per_interval`` - array of integer (int64) - - Yes - Array where each item corresponds to the total number of transpose elements processed by a group of instructions from the lower hierarchy level. Each corresponds to the time interval in the same position in the time_intervals field. Example: ``[256, 512, 128]``. * - ``hbm_transfer_size_per_interval`` - array of integer (int64) - bytes - Yes - Array where each item corresponds to the total number of HBM bytes transferred by a group of instructions from the lower hierarchy level. Each corresponds to the time interval in the same position in the time_intervals field. Example: ``[16384, 32768, 8192]``. * - ``hbm_read_bytes`` - integer (int64) - bytes - Yes - Total number of bytes read from HBM by all grouped instructions in this Framework instruction. * - ``hbm_write_bytes`` - integer (int64) - bytes - Yes - Total number of bytes written to HBM by all grouped instructions in this Framework instruction. * - ``sbuf_read_bytes`` - integer (int64) - bytes - No - Total number of bytes read from SBUF by all grouped instructions in this Framework instruction. * - ``sbuf_write_bytes`` - integer (int64) - bytes - No - Total number of bytes written to SBUF by all grouped instructions in this Framework instruction. * - ``load_weight_bytes`` - integer (int64) - bytes - No - Total number of bytes loaded by LDWEIGHTS instructions in this Framework instruction. * - ``tensor_engine_instruction_count`` - integer (int64) - - No - Total number of instructions executed by PE (Tensor) engine in this Framework instruction. * - ``scalar_engine_instruction_count`` - integer (int64) - - No - Total number of instructions executed by ACT (Scalar) engine in this Framework instruction. * - ``sync_engine_instruction_count`` - integer (int64) - - No - Total number of instructions executed by SP (Sync) engine in this Framework instruction. * - ``gpsimd_engine_instruction_count`` - integer (int64) - - No - Total number of instructions executed by POOL (GpSimd) engine in this Framework instruction. * - ``vector_engine_instruction_count`` - integer (int64) - - No - Total number of instructions executed by DVE (Vector) engine in this Framework instruction. .. _profile-schema-frameworknode: FrameworkNode ~~~~~~~~~~~~~ Framework hierarchy node information for displaying hierarchical structure of framework level operations. Each Node represents a / separated part of FrameworkInstruction. For example, if a FrameworkInstruction has the name layer[1]/layer[2], then layer[1] and layer[2] would be two FrameworkNodes. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``execution_index`` - integer (int32) - - Yes - Zero-based index of the NEFF execution this node belongs to. Matches execution_index in ExecutionInfo and NkiKernelCallSites. Use to scope framework nodes to a specific execution within a multi-execution profile. * - ``node_name`` - string - - Yes - Name of the current node in the hierarchy. Example: ``NeuronLlamaDecoderLayer``. * - ``children_names`` - string - - Yes - Comma separated array of child node names. Example: ``[child1, child2]``. * - ``parent_name`` - string - - No - Name of the parent node in the hierarchy. Example: ``parent_node``. * - ``children_ids`` - string - - Yes - Comma separated array of framework operator IDs associated with this node. Example: ``[123, 456, 789]``. * - ``start_ts`` - integer (int64) - nanoseconds - Yes - Earliest start time among all operations in this node. Unit is nanoseconds. * - ``end_ts`` - integer (int64) - nanoseconds - Yes - Latest end time among all operations in this node. Unit is nanoseconds. * - ``time_intervals`` - string - - Yes - A list of timestamp intervals after merging all corresponding instruction timestamps from the lower hierarchy level. Each interval is represented as [start, end] Example: ``[[1000, 2000], [3000, 4000]]``. * - ``element_count_per_interval`` - string - - No - Comma separated list where each item corresponds to the total number of elements processed by a group of instructions from the lower hierarchy level. Each corresponds to the time interval in the same position in the time_intervals field. Example: ``[1024, 2048, 512]``. * - ``transpose_element_count_per_interval`` - string - - No - Comma separated list where each item corresponds to the total number of transpose elements processed by a group of instructions from the lower hierarchy level. Each corresponds to the time interval in the same position in the time_intervals field. Example: ``[256, 512, 128]``. * - ``hbm_transfer_size_per_interval`` - string - bytes - No - Comma separated list where each item corresponds to the total number of HBM bytes transferred by a group of instructions from the lower hierarchy level. Each corresponds to the time interval in the same position in the time_intervals field. Example: ``[16384, 32768, 8192]``. * - ``hfu`` - number (float) - fraction (0-1) - No - Hardware FLOPs Utilization (HFU) percentage calculated once for a single operator, as (total element counts including transposes / (TENSOR_NUM_ROWS \* TENSOR_NUM_COLS \* total duration in cycles)). Range: 0 to 1. * - ``mfu`` - number (float) - fraction (0-1) - No - Model FLOPs Utilization (MFU) percentage calculated once for a single operator, as (total element counts excluding transposes / (TENSOR_NUM_ROWS \* TENSOR_NUM_COLS \* total duration in cycles)). Range: 0 to 1. * - ``hfu_per_interval`` - string - - No - Comma separated array of per-interval Hardware FLOPs Utilization (HFU) percentages corresponding to the time intervals in time_intervals field. Example: ``[23.5, 45.2, 67.8]``. * - ``mfu_per_interval`` - string - - No - Comma separated array of per-interval Model FLOPs Utilization (MFU) percentages corresponding to the time intervals in time_intervals field. Example: ``[20.1, 42.3, 65.4]``. * - ``mbu_per_interval`` - string - - No - Comma separated array of per-interval Memory Bandwidth Utilization (MBU) percentages corresponding to the time intervals in time_intervals field. Example: ``[15.2, 30.4, 45.6]``. * - ``hbm_read_bytes`` - integer (int64) - bytes - No - Total number of bytes read from HBM by all grouped instructions in this framework node. * - ``hbm_write_bytes`` - integer (int64) - bytes - No - Total number of bytes written to HBM by all grouped instructions in this framework node. * - ``sbuf_read_bytes`` - integer (int64) - bytes - No - Total number of bytes read from SBUF by all grouped instructions in this framework node. * - ``sbuf_write_bytes`` - integer (int64) - bytes - No - Total number of bytes written to SBUF by all grouped instructions in this framework node. * - ``tensor_engine_instruction_count`` - integer (int64) - - No - Total number of instructions executed by Tensor engine in this framework node. * - ``scalar_engine_instruction_count`` - integer (int64) - - No - Total number of instructions executed by Scalar engine in this framework node. * - ``sync_engine_instruction_count`` - integer (int64) - - No - Total number of instructions executed by Sync engine in this framework node. * - ``gpsimd_engine_instruction_count`` - integer (int64) - - No - Total number of instructions executed by GpSimd engine in this framework node. * - ``vector_engine_instruction_count`` - integer (int64) - - No - Total number of instructions executed by Vector engine in this framework node. * - ``total_time_tensor`` - string - nanoseconds - No - Comma separated array containing min and max timestamp for Tensor engine instructions in this framework node. Format is [min_timestamp, max_timestamp] or empty array [] if no Tensor instructions. Example: ``[1000000, 2000000]``. * - ``total_time_scalar`` - string - nanoseconds - No - Comma separated array containing min and max timestamp for Scalar engine instructions in this framework node. Format is [min_timestamp, max_timestamp] or empty array [] if no Scalar instructions. Example: ``[1500000, 2500000]``. * - ``total_time_sync`` - string - nanoseconds - No - Comma separated array containing min and max timestamp for Sync engine instructions in this framework node. Format is [min_timestamp, max_timestamp] or empty array [] if no Sync instructions. Example: ``[1200000, 1800000]``. * - ``total_time_gpsimd`` - string - nanoseconds - No - Comma separated array containing min and max timestamp for GpSimd engine instructions in this framework node. Format is [min_timestamp, max_timestamp] or empty array [] if no GpSimd instructions. Example: ``[1100000, 1900000]``. * - ``total_time_vector`` - string - nanoseconds - No - Comma separated array containing min and max timestamp for Vector engine instructions in this framework node. Format is [min_timestamp, max_timestamp] or empty array [] if no Vector instructions. Example: ``[1300000, 2100000]``. * - ``load_weight_bytes`` - integer (int64) - bytes - No - Total number of bytes loaded by LDWEIGHTS instructions in this framework node. * - ``total_channels_scalar`` - integer (int64) - - No - Total number of channels processed by Activation (Scalar) engine instructions in this framework node. * - ``total_channels_vector`` - integer (int64) - - No - Total number of channels processed by DVE (Vector) engine instructions in this framework node. .. _profile-schema-hbmallocation: HbmAllocation ~~~~~~~~~~~~~ HBM (High Bandwidth Memory) allocation periods showing which physical address ranges are occupied over time. Built from nrt_dma_mem_alloc and nrt_dma_mem_dealloc system trace events to track per-allocation occupancy. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``start_ts`` - integer (int64) - nanoseconds - Yes - Timestamp when the HBM address range becomes allocated. * - ``end_ts`` - integer (int64) - nanoseconds - Yes - Timestamp when the HBM address range becomes freed. Set to profile end timestamp if the allocation is never freed. * - ``start_address`` - integer (int64) - bytes - Yes - Starting physical memory address of the allocated region (inclusive - first byte of allocation). * - ``end_address`` - integer (int64) - bytes - Yes - Ending physical memory address of the allocated region (exclusive - start_address + size). * - ``usage_type`` - string - - Yes - Category of memory usage for this allocation. Examples include Tensors, Weights, Scratchpad, Code, Collectives, DMA Rings. Example: ``Tensors``. * - ``device_idx`` - integer (int64) - - Yes - Neuron device index (MLA index) that this allocation belongs to. * - ``hbm_idx`` - integer (int64) - - Yes - HBM channel index within the device. * - ``nc_idx`` - integer (int64) - - Yes - NeuronCore index within the device. * - ``tensor_name`` - string - - No - Name of the tensor or variable associated with this allocation (if available from debug info). Example: ``hidden_states``. * - ``model_id`` - integer (int64) - - No - Model ID that owns this allocation (if available). * - ``neff_name`` - string - - No - NEFF or operation name associated with this allocation (if available). Example: ``forward_pass``. * - ``pcore_idx`` - integer (int64) - - No - Index of the physical NeuronCore within a Logical NeuronCore (LNC). For LNC2 this field can have a value of either 0 or 1. * - ``trace_id`` - integer (int64) - - No - Systrace tracking_id correlation key. Matches the trace_id emitted in a PyTorch memory snapshot, linking this physical HBM allocation to a framework allocation. 0 when no correlation key was captured. * - ``source_frames`` - string - - No - Newline-joined source call stack (innermost first) for the allocation, sourced from a correlated PyTorch memory snapshot. Empty when no snapshot correlation exists. Example: ``rms_norm_neuron (rms_norm.py:18)``. * - ``virtual_address`` - integer (int64) - bytes - No - Framework-visible virtual address of the allocation from a correlated PyTorch memory snapshot. 0 when unavailable (runtime-internal allocations only have a physical address). .. _profile-schema-hbmusage: HbmUsage ~~~~~~~~ HBM time-series usage metrics such as throughput over time. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``timestamp`` - integer (int64) - nanoseconds - Yes - Timestamp for this measurement. * - ``current_throughput`` - number (float) - bytes_per_second - Yes - Rate of data transfer to/from HBM in bytes per second at this point of time. * - ``total_transfer_bytes`` - integer (int64) - bytes - Yes - Total bytes transferred to/from HBM, cumulative from the start of profiling to this timestamp. Each DMA transfer contributes the bytes it read when its source was in HBM plus the bytes it wrote when its destination was in HBM, so a transfer from HBM to HBM contributes both and one from SBUF to SBUF contributes neither. Summary.hbm_read_bytes and Summary.hbm_write_bytes count those same bytes, so the last point of this curve summed over pcore_idx equals Summary.hbm_read_bytes + Summary.hbm_write_bytes when every transfer's endpoints are known. * - ``pcore_idx`` - integer (int64) - - No - Index of the physical NeuronCore within a Logical NeuronCore (LNC). A Logical NeuronCore is a group of physical NeuronCores. For LNC2 this field can have a value of either 0 or 1. .. _profile-schema-hbmusagesummarybytype: HbmUsageSummaryByType ~~~~~~~~~~~~~~~~~~~~~ HBM memory usage information summarized by type. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``neuron_device_idx`` - integer (int64) - - Yes - The Neuron Device index. * - ``neuroncore_idx`` - integer (int64) - - Yes - The NeuronCore index. * - ``usage_type`` - string - - Yes - The type of memory usage. Example: ``Weights``. * - ``usage_bytes`` - integer (int64) - bytes - Yes - The amount of HBM used by this type of usage. .. _profile-schema-hloinstruction: HloInstruction ~~~~~~~~~~~~~~ HLO level instruction information for displaying hierarchy graph. Parent level is framework level, child level is tensorizer level. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``id`` - integer (int64) - - Yes - Unique identifier for corresponding assembly level instruction. * - ``raw_bir_id`` - integer (int64) - - Yes - Unique identifier for backend-level instruction. * - ``raw_penguin_id`` - integer (int64) - - Yes - Unique identifier for tensorizer-level instruction. * - ``raw_hlo_id`` - integer (int64) - - Yes - Unique identifier for hlo level instruction. * - ``raw_framework_id`` - integer (int64) - - Yes - Unique identifier for framework level instruction. * - ``hlo_name`` - string - - Yes - High-Level Optimizer (HLO) operation name. Example: ``dot.1``. * - ``modular_flow_invocation_idx`` - integer (int64) - - No - Integer for identifying which call is the current instruction. 0 represent first time calling the function, 1 represent second time calling the function. Example: ``1``. * - ``start_ts`` - integer (int64) - nanoseconds - Yes - Minimum start time among all of the corresponding instruction timestamps. Unit is nanoseconds. * - ``end_ts`` - integer (int64) - nanoseconds - Yes - Maximum end time among all of the corresponding instruction timestamps. Unit is nanoseconds. * - ``time_intervals`` - string - - Yes - A list of timestamp intervals after merging all corresponding instruction timestamps from the lower hierarchy level. Each interval is represented as [start, end]. Example: ``[[1, 10], [20, 50]]``. * - ``element_count_per_interval`` - array of integer (int64) - - Yes - Array where each item corresponds to the total number of elements processed by a group of instructions from the lower hierarchy level. Each corresponds to the time interval in the same position in the time_intervals field. Example: ``[1024, 2048, 512]``. * - ``transpose_element_count_per_interval`` - array of integer (int64) - - Yes - Array where each item corresponds to the total number of transpose elements processed by a group of instructions from the lower hierarchy level. Each corresponds to the time interval in the same position in the time_intervals field. Example: ``[256, 512, 128]``. * - ``hbm_transfer_size_per_interval`` - array of integer (int64) - bytes - Yes - Array where each item corresponds to the total number of HBM bytes transferred by a group of instructions from the lower hierarchy level. Each corresponds to the time interval in the same position in the time_intervals field. Example: ``[16384, 32768, 8192]``. * - ``hbm_read_bytes`` - integer (int64) - bytes - Yes - Total number of bytes read from HBM by all grouped instructions in this HLO instruction. * - ``hbm_write_bytes`` - integer (int64) - bytes - Yes - Total number of bytes written to HBM by all grouped instructions in this HLO instruction. * - ``sbuf_read_bytes`` - integer (int64) - bytes - No - Total number of bytes read from SBUF by all grouped instructions in this HLO instruction. * - ``sbuf_write_bytes`` - integer (int64) - bytes - No - Total number of bytes written to SBUF by all grouped instructions in this HLO instruction. * - ``load_weight_bytes`` - integer (int64) - bytes - No - Total number of bytes loaded by LDWEIGHTS instructions in this HLO instruction. .. _profile-schema-hostdevicedmausage: HostDeviceDmaUsage ~~~~~~~~~~~~~~~~~~ Host-device (top-level) DMA time-series usage metrics such as throughput over time. These DMAs move data between the host and the device as a whole and are not attributed to any single NeuronCore, so unlike DmaUsage there is no pcore_idx. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``timestamp`` - integer (int64) - nanoseconds - Yes - Timestamp for this measurement. * - ``current_throughput`` - number (float) - bytes_per_second - Yes - Rate of host-device data transfer in bytes per second at this timestamp. * - ``total_transfer_bytes`` - integer (int64) - bytes - Yes - Total bytes transferred over host-device DMA. This total is cumulative from the start of profiling to this timestamp. .. _profile-schema-hostmemusage: HostMemUsage ~~~~~~~~~~~~ Host memory utilization time-series metrics. Each row represents a memory usage sample for an instance. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``timestamp`` - integer (int64) - nanoseconds - Yes - Timestamp for this memory utilization measurement. * - ``used_bytes`` - integer (int64) - bytes - Yes - Amount of host memory currently in use. * - ``used_percent`` - number (float) - percent (0-100) - Yes - Host memory utilization percentage at this timestamp. Calculated as (used_bytes / memory_capacity) \* 100. Range: 0 to 100. * - ``instance_id`` - string - - Yes - Identifier for the instance on which this memory utilization was measured. Example: ``i-0b1ea78ca2865fd32``. * - ``ultraserver_id`` - integer (int64) - - No - Unique UltraServer ID if this measurement was captured on an UltraServer instance. * - ``ultraserver_instance_index`` - integer (int64) - - No - The instance index within the UltraServer on which this measurement was captured. .. _profile-schema-instruction: Instruction ~~~~~~~~~~~ Instructions executed on NeuronCore compute engines. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``execution_index`` - integer (int32) - - No - Zero-based index of the NEFF execution this instruction belongs to. Matches execution_index in ExecutionInfo, FrameworkNode, and NkiKernelCallSites. Use to scope instructions to a specific execution within a multi-execution profile. Null on profiles ingested before this field existed. * - ``start_ts`` - integer (int64) - nanoseconds - Yes - Timestamp when instruction started execution. * - ``end_ts`` - integer (int64) - nanoseconds - Yes - Timestamp when instruction finished execution. * - ``duration_ns`` - integer (int64) - nanoseconds - Yes - Duration of instruction execution. * - ``engine`` - string - - Yes - Compute engine that executed this instruction. Example: ``tensor``. * - ``timeline_track`` - string - - No - The name of the timeline track that this instruction appears on in the Neuron Explorer web UI. Example: ``TensorMatrix``. * - ``execution_unit`` - string (:ref:`ExecutionUnit `) - - No - Whether this instruction was executed by the sequencer (control plane) or dispatched to the compute engine (data plane). Sequencer instructions handle flow control, synchronization, and register manipulation. Compute instructions perform the actual data processing on the hardware engine. Allowed values: ``compute``, ``sequencer``. Example: ``compute``. * - ``sub_row`` - integer (int64) - - No - Logical partition used by timeline quantization. Instructions with different values are sampled independently within the same track, physical NeuronCore, and time bucket. Synthetic DGE phase spans use 1 so they do not compete with normal instructions; null is treated as 0. This field does not control UI row placement. * - ``pcore_idx`` - integer (int64) - - No - Index of the physical NeuronCore within a Logical NeuronCore (LNC). A Logical NeuronCore is a group of physical NeuronCores. For LNC2 this field can have a value of either 0 or 1. * - ``engine_instance_idx`` - integer (int64) - - Yes - The compute engine index that this instruction ran on. Default: ``0``. * - ``opcode`` - string - - Yes - Opcode for the instruction. Example: ``DRAIN``. * - ``operands`` - string - - Yes - Input operands for the instruction. Example: ``src0=bfloat16@0x2003000[1,1,1][1024,1,1] dst=bfloat16@0x3ef20[1,1,1][1024,1,1]``. * - ``symbolic_operands`` - string - - No - Contains the same value as the raw "operands" field, but with each device address replaced by the on-chip buffer that address refers to. Reading "operands" otherwise requires knowing the architecture's address map. Example: ``src0=bfloat16@PSUM(partition=0,bank=6,offset=0)[1,1,1][1024,1,1] dst=bfloat16@SBUF(partition=0)[1,1,1][1024,1,1]``. * - ``compiler_opcode`` - string - - No - Opcode of the instruction at compile time. The "compiler_opcode" and the runtime "opcode" can be different because certain missing or dynamic information is filled in at runtime. For example, the "compiler_opcode: PSEUDO_DMA_TRIGGER" may change to "opcode: WRITE" at runtime because that is when the write address is known. Example: ``PSEUDO_DMA_TRIGGER``. * - ``compiler_operands`` - string - - No - Operands of the instruction at compile time. The "compiler_operands" and the runtime "operands" can be different because certain missing or dynamic information is filled in at runtime. For example, the a compiler_opcode PSEUDO_DMA_TRIGGER with "comiler_operands: qSyncPIOParam0 block_id=1" may change to a opcode WRITE with "operands: data=0x1,0x1 size=4,4 dst=0xffff1009038,0xffff1029038" at runtime because that is when the write address is known. Example: ``name=sg0000 args_table_var_id=0``. * - ``id`` - integer (int64) - - Yes - Unique identifier for this instruction execution. * - ``pc`` - integer (int64) - - Yes - The runtime program counter is this instruction index on the given engine. * - ``compiler_pc`` - integer (int64) - - No - The compiler program counter is this instruction index on the given engine. The "compiler_pc" and the runtime "pc" can be different because the compiler can have instructions that are unrolled at runtime. Or the compiler can have a collective communication operation that is fulfilled by multiple runtime instructions. * - ``cci_slot`` - integer (int64) - - No - Slot of this compact control sub-instruction within the instruction at "compiler_pc". A compact control instruction packs several small control sub-instructions into one instruction word, and each sub-instruction executes as its own instruction with its own timing. Null for an instruction that is not a compact control sub-instruction. Example: ``12``. * - ``penguin_id`` - string - - No - Tensorizer-level compiler ID containing both raw ID and context information. Example: ``sg0000:83``. * - ``raw_penguin_id`` - integer (int64) - - No - Raw internal tensorizer-level compiler identifier. * - ``bir_id`` - string - - No - Backend-level instruction identifier. Example: ``sg0000:218805``. * - ``raw_bir_id`` - integer (int64) - - No - Unprocessed backend-level instruction ID containing both raw ID and context information. This is a 64-bit integer stored as a string. Bits 63:48 encode the subgraph ID. Bits 47:32 encode the function ID. Bits 31:0 encode the instruction ID within the function. * - ``bir_instruction_name`` - string - - No - Name of the backend-level instruction, taken from the compiler debug info. Example: ``I-83-1_sg0001``. * - ``kernel_instruction_name`` - string - - No - Unique kernel instruction identifier formatted as ``bir_instruction_name:parent_tensorizer_id``, where parent_tensorizer_id is the parent tensorizer-level instruction ID for this instruction. Corresponds to the ``kernel_instruction_name`` field in the KernelInstructions table. Example: ``inst.42:67``. * - ``nki_call_site_id`` - integer (int64) - - No - ID of the NKI kernel call site this instruction belongs to. Null if this instruction is not part of an NKI kernel. Corresponds to the ``nki_call_site_id`` field in the NkiKernelCallSites table. * - ``crc_in`` - integer (int64) - - No - In CRC for integrity checking. It is a 16-bit hash value for the input data of this instruction. The crc_in value should match between runs for the same exact instruction call. * - ``crc_out`` - integer (int64) - - No - Out CRC for integrity verification. It is a 16-bit hash value for the output data of this instruction. The crc_out value should match between runs for the same exact instruction call. * - ``evt_wait_time_ns`` - integer (int64) - nanoseconds - No - Time spent waiting for events or synchronization. The duration between the time the instruction is ready to be pushed into the engine execution queue and the time the event/semaphore condition is met. * - ``hlo_attrs`` - string - - No - The debug info associated with the HLO instruction that produced this instruction. Example: ``transpose=true,axis=[0,2,1]``. * - ``hlo_name`` - string - - No - High-Level Optimizer (HLO) operation name. Example: ``dot.1``. * - ``raw_hlo_id`` - integer (int64) - - No - Unique identifier for hlo level instruction. * - ``raw_framework_id`` - integer (int64) - - No - Unique identifier for framework level instruction. When an instruction maps to multiple framework IDs, the ID is selected using modular flow invocation index. * - ``layer`` - string - - No - Here "layer" means the name from the framework level debug info. Example: ``layer_1``. * - ``stack_frame_ids`` - string - - No - List of stack frame identifiers for debugging. An instruction can map back to multiple stack frame ids due to compiler optimizations such as operator fusion. Example: ``[101, 102, 103]``. * - ``nki_source_location`` - string - - No - Neuron Kernel Interface (NKI) source code location that generated this instruction. This is a list of strings because an instruction can map back to multiple source locations due to compiler optimizations such as operator fusion. Example: ``file.py:42``. * - ``bir_debug_info_source_location`` - string - - No - Source code location recorded in compiler debug info. This field can reference various source file types including NKI files, non-NKI Python files, and C++ files. Example: ``file.py:42``. * - ``dma_trigger_start_ts`` - integer (int64) - - No - Start timestamp of the DMA that this instruction triggered. * - ``dma_last_trigger_start_ts`` - integer (int64) - - No - Start timestamp of the last DMA triggered by this instruction. * - ``dma_wait_time_ns`` - integer (int64) - nanoseconds - No - The amount of time delay between the end timestamp of the trigger compute engine instruction to the start timestamp of the DMA that was triggered. * - ``instruction_chain`` - string - - No - List of instructions in this chain. Some instructions must be executed together in a sequential order such as custom operations or pseuodo operations. This field lists all instructions in the executed chain in the following format "compiler_opcode compiler_operands\\n". Newline is used to separate instructions in the chain. Example: ``PSEUDO_TRIGGER_COLLECTIVE_2 $S[15]>=1 $S[15]++@complete ctype=ALL_REDUCE op=ADD dtype=bfloat16 cc_dim=0 src0_var_id=1353 src0_var_offset=0 src0_shape_reg=0 src0_pattern=[1,1][8192,1]\n PSEUDO_EXTENSION group_id=0 channel_id=0 stream_id=0 dst_var_id=1354 dst_var_offset=0 dst_shape_reg=0 dst_pattern=[1,1][8192,1]``. * - ``hbm_read_bytes`` - integer (int64) - - No - Number of bytes read from HBM by this instruction. * - ``hbm_write_bytes`` - integer (int64) - - No - Number of bytes written to HBM by this instruction. * - ``sbuf_read_bytes`` - integer (int64) - - No - Number of bytes read from SBUF by this instruction. * - ``sbuf_write_bytes`` - integer (int64) - - No - Number of bytes written to SBUF by this instruction. * - ``spill_save_bytes`` - integer (int64) - - No - Number of bytes saved to HBM from SBUF by this instruction due to spill (lack of space in SBUF). * - ``spill_reload_bytes`` - integer (int64) - - No - Number of bytes loaded from HBM to SBUF by this instruction due to spill (lack of space in SBUF). * - ``psum_read_bytes`` - integer (int64) - - No - Number of bytes read from PSUM by this instruction. * - ``psum_write_bytes`` - integer (int64) - - No - Number of bytes written to PSUM by this instruction. * - ``psum_read_sbuf_write_bytes`` - integer (int64) - - No - Number of bytes read from PSUM into SBUF by this instruction. * - ``sbuf_read_partition_start`` - integer (int32) - - No - Lowest SBUF partition index read by this instruction. A tile can only start at partition 0, 32, 64, or 96. Null when no SBUF read address could be decoded. Null on performance-simulation profiles. * - ``sbuf_read_partition_end`` - integer (int32) - - No - Highest SBUF partition index read by this instruction, inclusive. Equal to sbuf_read_partition_start for a single-partition read. Null when sbuf_read_partition_start is null. * - ``sbuf_write_partition_start`` - integer (int32) - - No - Lowest SBUF partition index written by this instruction. A tile can only start at partition 0, 32, 64, or 96. Null when no SBUF write address could be decoded. Null on performance-simulation profiles. * - ``sbuf_write_partition_end`` - integer (int32) - - No - Highest SBUF partition index written by this instruction, inclusive. Equal to sbuf_write_partition_start for a single-partition write. Null when sbuf_write_partition_start is null. * - ``psum_read_partition_start`` - integer (int32) - - No - Lowest PSUM partition index read by this instruction. A tile can only start at partition 0, 32, 64, or 96. Null when no PSUM read address could be decoded. Null on performance-simulation profiles. * - ``psum_read_partition_end`` - integer (int32) - - No - Highest PSUM partition index read by this instruction, inclusive. Equal to psum_read_partition_start for a single-partition read. Null when psum_read_partition_start is null. * - ``psum_write_partition_start`` - integer (int32) - - No - Lowest PSUM partition index written by this instruction. A tile can only start at partition 0, 32, 64, or 96. Null when no PSUM write address could be decoded. Null on performance-simulation profiles. * - ``psum_write_partition_end`` - integer (int32) - - No - Highest PSUM partition index written by this instruction, inclusive. Equal to psum_write_partition_start for a single-partition write. Null when psum_write_partition_start is null. * - ``psum_read_banks`` - array of integer (int64) - - No - Every PSUM bank this instruction reads, ascending, across all of its read operands. Empty when no PSUM read address could be decoded, and on performance-simulation profiles. Requires the operand's element width, so it is empty for a read operand whose data type is not printed even when the partition range is known. Example: ``[5]``. * - ``psum_write_banks`` - array of integer (int64) - - No - Every PSUM bank this instruction writes, ascending. Empty when no PSUM write address could be decoded, and on performance-simulation profiles. Example: ``[2]``. * - ``input_queue_bytes`` - integer (int64) - - No - Number of bytes loaded from an INPUT DMA queue by this instruction. * - ``output_queue_bytes`` - integer (int64) - - No - Number of bytes loaded from an OUTPUT DMA queue by this instruction. * - ``weight_queue_bytes`` - integer (int64) - - No - Number of bytes loaded from a DATA / WEIGHT DMA queue by this instruction. * - ``performance_mode`` - string (:ref:`PerformanceMode `) - - No - Performance mode for this instruction. Possible values are "2x" (double throughput), "4x" (quad throughput), or null (base mode / not applicable). Only applies to Vector and Scalar Engine instructions. Allowed values: ``2x``, ``4x``. * - ``tensor_instruction_type`` - string - - No - Instruction type name from the debug info. Example: ``["REGULAR", "TRANSPOSE"]``. * - ``tensor_elements`` - integer (int64) - - No - Do not use this field, it will eventually be deprecated as it can be misleading depending on data type and architecture, instead use raw_flops and adjusted_flops. Number of data elements processed by this instruction. Can be used to get total size of a MATMUL operation. For example a MATMUL with vector sizes 128 \* 128 has 16384 elements. * - ``raw_flops`` - integer (int64) - - No - Number of floating-point operations (FLOPs) performed by this instruction. Each MAC (multiply-accumulate) operation counts as 2 FLOPs (one multiply, one add). Note that as each FP32 matmul lowers to two FP32_LO_HI instructions, each FP32_LO_HI matmul records half the logical FLOPs. * - ``adjusted_flops`` - integer (int64) - - No - BF16-equivalent FLOPs for this instruction. Scales FLOPs to the BF16 throughput equivalent so the result can be used with BF16 peak throughput as the denominator for PE utilization calculations. BF16, FP16, and TF32 use a 1.0x multiplier. FP32 uses 4.0x. On Trn2/Trn3, FP8@FP8 and UINT8@UINT8 matmuls use 0.5x; FP8 or UINT8 mixed with a higher-precision operand uses 1.0x (executes at BF16 throughput). FP8 uses 1.0x on Trn1/Inf2. MXFP8 and MXFP4 use 0.25x on Trn3. * - ``tensor_elements_transpose`` - integer (int64) - - No - Number of data elements processed by this instruction, only set when tensor_instruction_type is TRANSPOSE. 0 for non-transpose instructions. * - ``flops_transpose`` - integer (int64) - - No - Number of transpose floating-point operations (FLOPs) performed by this instruction, only set when tensor_instruction_type is TRANSPOSE. 0 for non-transpose instructions. * - ``adjusted_flops_transpose`` - integer (int64) - - No - BF16-equivalent transpose FLOPs for this instruction, only set when tensor_instruction_type is TRANSPOSE and 0 otherwise. Uses the same scaling multipliers as adjusted_flops. * - ``partition_dimension`` - integer (int64) - - No - Partition-axis extent of this instruction's tensor access, decoded from the instruction's ISA representation. For MATMUL/LDWEIGHTS this is the number of active PE rows; for channel-oriented ops (activation, tensor-tensor, pooling) it is the number of active channels. Null for instructions whose partition dimension is not applicable or cannot be decoded. * - ``free_dimension`` - integer (int64) - - No - Free-axis extent of this instruction's tensor access, decoded from the instruction's ISA representation. For MATMUL/LDWEIGHTS this is the free-axis extent of the source memory pattern (the moving matrix for MATMUL, the stationary weights for LDWEIGHTS), i.e. the product of its non-partition axis extents; for other ops it is the product of the destination's non-partition (free) axis extents. Null for instructions whose free dimension is not applicable or cannot be decoded. * - ``src1_dtype`` - string (:ref:`DataType `) - - No - Data type of this instruction's first input operand, decoded from the instruction's ISA representation. For MATMUL this is the stationary (weights) operand. Unset for instructions whose input data type is not applicable or cannot be decoded. Allowed values: ``fp64``, ``fp32``, ``fp32r``, ``tf32``, ``fp16``, ``bf16``, ``fp8e3``, ``fp8e4``, ``fp8e5``, ``mxfp8``, ``mxfp4``, ``int8``, ``int16``, ``int32``, ``int64``, ``uint8``, ``uint16``, ``uint32``, ``uint64``, ``unknown``. Example: ``bf16``. * - ``src2_dtype`` - string (:ref:`DataType `) - - No - Data type of this instruction's second input operand, decoded from the instruction's ISA representation. For MATMUL this is the moving (activation) operand; for tensor-tensor ops it is the second input tensor. Unset for single-input instructions or when it cannot be decoded. Allowed values: ``fp64``, ``fp32``, ``fp32r``, ``tf32``, ``fp16``, ``bf16``, ``fp8e3``, ``fp8e4``, ``fp8e5``, ``mxfp8``, ``mxfp4``, ``int8``, ``int16``, ``int32``, ``int64``, ``uint8``, ``uint16``, ``uint32``, ``uint64``, ``unknown``. Example: ``bf16``. * - ``dst_dtype`` - string (:ref:`DataType `) - - No - Data type of this instruction's output operand, decoded from the instruction's ISA representation. Unset for instructions whose output data type is not applicable or cannot be decoded. Allowed values: ``fp64``, ``fp32``, ``fp32r``, ``tf32``, ``fp16``, ``bf16``, ``fp8e3``, ``fp8e4``, ``fp8e5``, ``mxfp8``, ``mxfp4``, ``int8``, ``int16``, ``int32``, ``int64``, ``uint8``, ``uint16``, ``uint32``, ``uint64``, ``unknown``. Example: ``fp32``. * - ``scalar_activation_fn`` - string - - No - Type of activation function used. Example: ``COPY``. .. _profile-schema-kernelinstructions: KernelInstructions ~~~~~~~~~~~~~~~~~~ Kernel instructions with references to their call stack and iteration context. Each row represents one instruction that references stack frames and iteration variables defining its execution context. Example: Instruction "inst.72:0" located at line 76 of mlp_with_mm_kernel.py, with stack_frame_id=42 (linking to KernelStackFrames table), iteration_variables_id=156 (linking to KernelIterationVariables table), stack_frame_depth=1 (single function in call stack), iteration_variable_depth=2 (inside two nested loops). .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``kernel_instruction_name`` - string - - Yes - Unique kernel instruction identifier formatted as ``instruction_name:parent_tensorizer_id``, where parent_tensorizer_id is the parent tensorizer-level instruction ID for this instruction. parent_tensorizer_id refers to the parent. Example: ``inst.72:0``. * - ``stack_frame_id`` - integer (int64) - - No - Reference to the deepest stack frame in the call stack for this instruction. Links to the KernelStackFrames table. Null if no stack frame information is available. Example: ``42``. * - ``stack_frame_depth`` - integer (int64) - - No - Total depth of the call stack (number of frames in the stack). Example: ``1``. * - ``iteration_variables_id`` - integer (int64) - - No - Reference to the innermost iteration variable in the nested loop structure for this instruction. Links to the KernelIterationVariables table. Null if no iteration variable information is available. Example: ``156``. * - ``iteration_variable_depth`` - integer (int64) - - No - Total depth of iteration variable nesting (number of nested loops). Example: ``2``. * - ``pcore_idx`` - integer (int64) - - No - Index of the physical NeuronCore within a Logical NeuronCore (LNC). A Logical NeuronCore is a group of physical NeuronCores. For LNC2 this field can have a value of either 0 or 1. * - ``file_path`` - string - - No - Source file path where the code that generates this instruction is located. Example: ``/home/ubuntu/mlp_with_mm_kernel.py``. * - ``line_number`` - integer (int64) - - No - Line number in the source file that generates this instruction. Example: ``76``. * - ``annotation_labels`` - array of string - - No - NKI Region labels active when this instruction was emitted. Set by ``nki.language.profiler_region_open`` and ``nki.language.profiler_region_close`` during kernel compilation. Empty or absent when the instruction is not inside any region. Example: ``["compute", "per_iter"]``. * - ``annotation_colors`` - array of string - - No - User-supplied colors for the regions, corresponding 1:1 with ``annotation_labels``. Each color is a hex string, or empty if no color was specified. Example: ``["#A133FF", "#FF5733"]``. * - ``annotation_unbalanced`` - array of string - - No - Region labels on this instruction that are unbalanced, meaning a ``profiler_region_open`` without a matching ``profiler_region_close``. A subset of the ``annotation_labels`` base names. Empty or absent when all regions are balanced. Example: ``["unbalanced_bad"]``. .. _profile-schema-kerneliterationvariables: KernelIterationVariables ~~~~~~~~~~~~~~~~~~~~~~~~ Loop iteration variables for kernel instructions forming a loop nest. Each row represents one loop variable's value at a specific iteration. When loops are nested, parent_iteration_variable_id links to the outer loop variable. Example: Loop variable "k" defined at line 72 of mlp_with_mm_kernel.py has value 3, with iteration_variable_id=102 and parent_iteration_variable_id=101 (entry of the outer loop variable). .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``iteration_variable_id`` - integer (int64) - - Yes - Unique identifier for this iteration variable entry. Example: ``102``. * - ``variable_name`` - string - - No - Name of the iteration variable. Example: ``k``. * - ``variable_value`` - integer (int64) - - No - Value of the iteration variable for this specific iteration. Example: ``3``. * - ``parent_iteration_variable_id`` - integer (int64) - - No - Reference to the outer loop's iteration variable. Null for the outermost loop. Example: ``101``. * - ``pcore_idx`` - integer (int64) - - No - Index of the physical NeuronCore within a Logical NeuronCore (LNC). A Logical NeuronCore is a group of physical NeuronCores. For LNC2 this field can have a value of either 0 or 1. * - ``file_path`` - string - - No - Source file path where this iteration variable is defined. Example: ``/home/ubuntu/mlp_with_mm_kernel.py``. * - ``line_number`` - integer (int64) - - No - Line number in the source file where this iteration variable is defined. Example: ``72``. .. _profile-schema-kernelstackframes: KernelStackFrames ~~~~~~~~~~~~~~~~~ Call stack frames for kernel instructions forming a call stack. Each row represents one function in the call stack. When functions are nested, parent_stack_frame_id links to the caller function. Example: Function "output_projection_cte" defined at line 17 of mlp_with_mm_kernel.py, with stack_frame_id=42 and parent_stack_frame_id=41 (entry of the caller function). .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``stack_frame_id`` - integer (int64) - - Yes - Unique identifier for this stack frame entry. Example: ``42``. * - ``fully_qualified_function_name`` - string - - No - Fully qualified function name for this stack frame. Example: ``nkilib.experimental.conv.conv3d.conv3d``. * - ``parent_stack_frame_id`` - integer (int64) - - No - Reference to the parent (caller) stack frame in the call stack. Null for the outermost function. Example: ``41``. * - ``pcore_idx`` - integer (int64) - - No - Index of the physical NeuronCore within a Logical NeuronCore (LNC). A Logical NeuronCore is a group of physical NeuronCores. For LNC2 this field can have a value of either 0 or 1. * - ``file_path`` - string - - No - Source file path where this function is defined. Example: ``/home/ubuntu/mlp_with_mm_kernel.py``. * - ``function_line_number`` - integer (int64) - - No - The line number of the definition line of this frame's function. Not to be confused with the parent stack frame's ``line_number``, which is one level up in the call stack. Example: ``17``. * - ``line_number`` - integer (int64) - - No - The line where this function frame of the stack was called, so two calls to the same function from different lines are distinguishable. Null for the outermost function, which nothing called. Example: ``124``. .. _profile-schema-memorybandwidthpoint: MemoryBandwidthPoint ~~~~~~~~~~~~~~~~~~~~ A single sample in a memory bandwidth time-series. Represents the aggregate bandwidth of all active tensor I/O transfers at a given timestamp. This is a transient type computed at query time from system profile dmem_buf_copyin/dmem_buf_copyout events — it is not persisted to a parquet table. The server performs the aggregation so the client receives ready-to-plot data. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``timestamp`` - integer (int64) - nanoseconds - Yes - Timestamp for this bandwidth measurement. Matches system profile trace event timestamps. * - ``bandwidth_bytes_per_sec`` - number (double) - bytes per second - Yes - Aggregate bandwidth at this timestamp, computed as the sum of bandwidths of all tensor I/O transfers active at this point in time. .. _profile-schema-memorybandwidthresponse: MemoryBandwidthResponse ~~~~~~~~~~~~~~~~~~~~~~~ Response payload for the memory bandwidth API endpoint. Contains time-series data for host-to-device, device-to-host, and bidirectional memory bandwidth computed from system profile tensor I/O events (dmem_buf_copyin / dmem_buf_copyout). This is a transient, computed response — it is not persisted to a parquet table. The server aggregates raw events into grouped, sorted time-series so the client can render bandwidth plots directly without additional processing. The grouping level (per neuron core, per instance, etc.) is controlled by the group_by query parameter. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``series`` - array of :ref:`MemoryBandwidthSeries ` - - Yes - List of bandwidth time-series. Each series represents a unique combination of grouping key and transfer direction. Example: ``[{"instance_id": "i-0b1ea78ca2865fd32", "worker_gid": 0, "nc_idx": 0, "direction": "host_to_device", "points": [{"timestamp": 1000, "bandwidth_bytes_per_sec": "2.048e9"}]}]``. .. _profile-schema-memorybandwidthseries: MemoryBandwidthSeries ~~~~~~~~~~~~~~~~~~~~~ A time-ordered sequence of bandwidth samples for a single grouping of system profile tensor I/O events. The grouping is controlled by the group_by query parameter and is computed server-side so the client receives pre-aggregated, ready-to-plot data without additional processing. Current grouping modes produce one series per (instance_id, worker_gid, direction) or per (instance_id, direction). Future modes may support per-device, per-cluster, or other aggregation levels. This is a transient type — it is not persisted to a parquet table. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``instance_id`` - string - - Yes - Identifier for the instance whose events are aggregated in this series. Example: ``i-0b1ea78ca2865fd32``. * - ``worker_gid`` - integer (int64) - - No - Global worker identifier the events belong to. Unique across all processes and instances, so it is the correct grouping key for per-core bandwidth in multi-process workloads. Absent when the grouping aggregates across workers (e.g. group_by=instance) or for profiles captured before worker_gid was populated. * - ``nc_idx`` - integer (int64) - - No - Logical NeuronCore index (LNC) the events occurred on. Process-local — not unique across an instance when multiple processes run. Provided for display alongside worker_gid. Absent when the grouping aggregates across cores (e.g. group_by=instance). * - ``direction`` - :ref:`MemoryBandwidthDirection ` - - Yes - Direction of a memory bandwidth transfer between host and device. 'host_to_device' corresponds to dmem_buf_copyin events (system profile). 'device_to_host' corresponds to dmem_buf_copyout events (system profile). 'bidirectional' is the simultaneous sum of both directions and only appears as a computed series direction, never on an individual event. Allowed values: ``host_to_device``, ``device_to_host``, ``bidirectional``. * - ``points`` - array of :ref:`MemoryBandwidthPoint ` - - Yes - Time-ordered bandwidth samples. Each point represents the aggregate bandwidth at a unique timestamp where an event starts or ends. Example: ``[{"timestamp": 1766072156592459482, "bandwidth_bytes_per_sec": "2.048e9"}, {"timestamp": 1766072156692459482, "bandwidth_bytes_per_sec": 0}]``. .. _profile-schema-metadata: Metadata ~~~~~~~~ Metadata associated with this profile, including details about the system configuration. Also see the NeffHeader table for metadata included in the NEFF header. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``profile_schema_version`` - string - - Yes - The version of the data schema, which is the layout of the data contained within this profile. See the table "SchemaFields.parquet" for a full listing of what data tables and fields exist in this schema version. Different values may have a different schema. Versioning scheme: MAJOR.MINOR.PATCH MAJOR: incompatible change - table removal or rename, - field removal or rename, - field type change, - field becomes optional. MINOR: compatible addition - new table, - new field, - field becomes required. PATCH: compatible change - any field metadata change (description, example, minimum, unit, etc.), - parquet compression, encoding, etc. (doesn't change data values, just how data is stored), - spacing, formatting, or ordering change to schema definition yaml schema files. Also note, things that this schema version does not track, but is tracked by the "profiler_version" field: - the actual data values that are produced, - the implementation that produced the data, - the performance of producing the data. Example: ``1.0.0``. * - ``ntff_version`` - integer (int32) - - Yes - Version of the NTFF format used. This is the version of the raw profile data before processing. * - ``num_physical_cores`` - integer (int32) - - Yes - The number of physical NeuronCores in the Logical NeuronCore. * - ``vnc_enabled`` - boolean - - Yes - Whether execution was on Logical NeuronCores consisting of multiple physical cores. Example: ``true``. * - ``first_hw_timestamp`` - integer (int64) - nanoseconds - Yes - First timestamp in the trace. * - ``last_hw_timestamp`` - integer (int64) - nanoseconds - Yes - Last timestamp in the trace. * - ``first_ts`` - string - - Yes - First timestamp in Unix time format. Example: ``1970-01-01 00:00:00+00:00``. * - ``last_ts`` - string - - Yes - Last timestamp in Unix time format. Example: ``1970-01-01 00:00:00+00:00``. * - ``dma_ddr_bandwidth`` - number (float) - bytes_per_second - Yes - Maximum DMA DDR bandwidth available to a single physical NeuronCore on the system on which the profile was captured. * - ``num_dma_engines_per_neuroncore`` - integer (int32) - - Yes - Number of DMA engines per physical NeuronCore. * - ``num_engine_instances`` - integer (int32) - - No - Largest number of instances of a single compute engine type within a physical NeuronCore, as declared by the compiler for the NEFFs in this profile. Minimum: 1. Example: ``4``. * - ``hbm_ddr_bandwidth`` - number (float) - bytes_per_second - Yes - Maximum HBM bandwidth on the system on which the profile was captured. This is for a single HBM stack. * - ``tensor_engine_num_rows`` - integer (int64) - - Yes - The number of rows in the systolic array of the Tensor Engine. * - ``tensor_engine_num_cols`` - integer (int64) - - Yes - The number of columns in the systolic array of the Tensor Engine. * - ``tensor_engine_clock_freq`` - number (float) - gigahertz - Yes - The clock frequency of the Tensor Engine. * - ``vector_engine_clock_freq`` - number (float) - gigahertz - Yes - The clock frequency of the Vector Engine. * - ``vector_engine_num_channels`` - integer (int64) - - Yes - The number of channels in the Vector Engine. * - ``scalar_engine_clock_freq`` - number (float) - gigahertz - Yes - The clock frequency of the Scalar Engine. * - ``scalar_engine_num_channels`` - integer (int64) - - Yes - The number of channels in the Scalar Engine. * - ``gpsimd_engine_clock_freq`` - number (float) - gigahertz - Yes - The clock frequency of the GpSimd Engine. * - ``gpsimd_engine_num_channels`` - integer (int64) - - Yes - The number of channels in the GpSimd Engine. * - ``neuroncore_clock_freq`` - number (float) - gigahertz - Yes - The clock frequency of the NeuronCore. * - ``ticks_per_nanosec`` - integer (int64) - - Yes - The number of ticks per nanosecond that the profile captures. Used to convert internal profile timestamps to nanoseconds. * - ``instance_type`` - string - - No - The instance type on which this profile was captured. Example: ``trn1.2xlarge``. * - ``neuroncore_version`` - string - - No - The NeuronCore architecture version of the hardware on which this profile was captured, for example v2 for trn1 and inf2, or v3 for trn2. Independent of the Logical NeuronCore configuration. Prefer this over instance_type when keying behavior on the architecture. Example: ``v3``. * - ``runtime_version`` - string - - No - The Neuron Runtime library version with which this profile was captured. Example: ``2.19.1``. * - ``collectives_version`` - string - - No - The Neuron Collectives library version with which this profile was captured. Example: ``2.20.22``. * - ``driver_version`` - string - - No - The Neuron driver version with which this profile was captured. Example: ``2.15.9``. * - ``explorer_version`` - string - - No - The Neuron Explorer version which post-processed this captured profile. Example: ``2.19.1``. * - ``compiler_version`` - string - - No - The Neuron Compiler version which built the NEFF that was executed during this profile. This is the tool_version field in the NEFF's info.json file. Example: ``2.13.2``. * - ``nki_version`` - string - - No - The Neuron Kernel Interface (NKI) version used to author the kernels executed during this profile. Aids debugging and reproducibility by identifying which NKI version generated the profile. Optional; only present when the profile carries NKI version metadata. Example: ``0.6.0``. * - ``neuroncore_idx`` - integer (int64) - - No - The index of this NeuronCore. * - ``neuron_device_idx`` - integer (int64) - - No - The index of this NeuronDevice. * - ``ultraserver_id`` - integer (int64) - - No - Unique UltraServer ID if this profile was captured on an UltraServer instance. * - ``ultraserver_type`` - string - - No - The type of instances connectivity within an UltraServer. Example: ``UltraCluster``. * - ``ultraserver_size`` - integer (int32) - - No - The total numbers of instance in an UltraServer. * - ``ultraserver_instance_index`` - integer (int32) - - No - The instance index within the UltraServer on which this profile was captured. * - ``s3_bucket_path`` - string - - No - Full S3 path where profile files (NEFF/NTFF/source/debug) are stored and uploaded. Example: ``s3://my-bucket/profile/``. * - ``profile_name`` - string - - No - User chosen name for this profile. Example: ``mlp_nki_mm``. * - ``explorer_build_time`` - string - - No - The date and time when the ``neuron-explorer`` tool was built. Example: ``2025-11-14 15:14:53+00:00``. * - ``capture_type`` - string (:ref:`CaptureType `) - - No - Whether the profiled activity ran on a Neuron device or in a simulator. Allowed values: ``hardware``, ``s0``, ``perf-sim``. Example: ``hardware``. .. _profile-schema-neffheader: NeffHeader ~~~~~~~~~~ Information contained in the NEFF header. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``network_name`` - string - - No - Name of the NEFF. Example: ``file.neff``. * - ``uuid`` - string - - No - Unique identifier of the NEFF. Example: ``ae5a8d94ac9c4104b0e01ece540d4097``. * - ``hash`` - string - - No - Hash of the NEFF. Example: ``6879b275cffd23c20fa431a0b1472002``. * - ``enabled_features`` - string - - No - Comma-separated list of enabled features. Example: ``coalesced-cc, queue-set-instances, on-chip-functions``. * - ``header_size`` - integer (int64) - bytes - No - Size of the NEFF header in bytes. * - ``data_size`` - integer (int64) - bytes - No - Data size in bytes. * - ``major_version`` - integer (int64) - - No - NEFF major version number. * - ``minor_version`` - integer (int64) - - No - NEFF minor version number. * - ``package_version`` - string - - No - Package version. Example: ``0``. * - ``build_version`` - string - - No - Build version. Example: ``0``. * - ``cc_participant_count`` - integer (int64) - - No - Number of collective communication (CC) participants involved. * - ``neuroncore_count`` - integer (int64) - - No - Number of NeuronCores used by the NEFF. * - ``logical_neuroncore_count`` - integer (int64) - - No - Number of Logical NeuronCores used by the NEFF. * - ``pcore_count_per_lnc`` - integer (int64) - - No - Number of physical NeuronCores per Logical NeuronCore (LNC). .. _profile-schema-nkikernelcallsites: NkiKernelCallSites ~~~~~~~~~~~~~~~~~~ One row per NKI kernel call site per NEFF execution. Each row represents one place in the compiled graph where an NKI kernel was called. Two calls to the same kernel function within an execution produce two rows with different nki_call_site_id values. Use execution_index to scope rows to a specific NEFF execution within the profile. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``execution_index`` - integer (int32) - - Yes - Zero-based index of the execution this call site belongs to. Matches execution_index in ExecutionInfo. Use with nki_call_site_id to uniquely identify a call site across a multi-NEFF profile. * - ``neff_name`` - string - - No - File name of the NEFF that contains this call site. Matches neff_name in ExecutionInfo. Example: ``model.neff``. * - ``nki_call_site_id`` - integer (int64) - - Yes - ID of the NKI kernel call site this instruction belongs to. Unique within a single NEFF execution; use with execution_index to scope across executions. * - ``kernel_name`` - string - - No - Name of the NKI kernel. Null when kernel_debug_info.json is absent from the NEFF. Example: ``nki_matmul``. * - ``parent_framework_ids`` - array of string - - No - Raw framework IDs of the parent framework-level nodes whose instructions this kernel ran within. Each ID corresponds to a raw_framework_id in the FrameworkNode table. Empty when the profile has no framework-level debug info. Example: ``["42", "57"]``. * - ``start_ts`` - integer (int64) - nanoseconds - No - Earliest start time among all instructions belonging to this call site. Unit is nanoseconds. Null when the call site has no timed instructions. * - ``end_ts`` - integer (int64) - nanoseconds - No - Latest end time among all instructions belonging to this call site. Unit is nanoseconds. Null when the call site has no timed instructions. .. _profile-schema-opcodesummary: OpcodeSummary ~~~~~~~~~~~~~ Summary of instruction counts and total execution time grouped by compute engine and opcode. Each row represents one (engine, opcode) combination with the total number of instructions and their cumulative execution time. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``engine`` - string - - Yes - Compute engine that executed the instructions. Example: ``tensor``. * - ``opcode`` - string - - Yes - Opcode of the instructions. Example: ``MATMUL``. * - ``instruction_count`` - integer (int64) - - Yes - Total number of instructions with this engine and opcode combination. * - ``instruction_time_ns`` - integer (int64) - nanoseconds - Yes - Total cumulative execution time of all instructions with this engine and opcode combination. Overlapping instructions on the same engine are added together (not merged). * - ``active_time_ns`` - integer (int64) - nanoseconds - Yes - Duration of time when at least one instruction with this engine and opcode combination is executing. Overlapping instructions are merged so that concurrent execution is counted only once. * - ``active_time_percent`` - number (float) - - Yes - Percentage of total execution time when at least one instruction with this engine and opcode combination is executing. Range: 0 to 1. .. _profile-schema-pendingdma: PendingDma ~~~~~~~~~~ Outstanding DMA operations that are queued. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``timestamp`` - integer (int64) - nanoseconds - Yes - Time of pending DMA measurement. * - ``value`` - integer (int32) - - Yes - Number of pending DMA operations queued. * - ``pcore_idx`` - integer (int64) - - No - Index of the physical NeuronCore within a Logical NeuronCore (LNC). A Logical NeuronCore is a group of physical NeuronCores. For LNC2 this field can have a value of either 0 or 1. .. _profile-schema-penguininstruction: PenguinInstruction ~~~~~~~~~~~~~~~~~~ Tensorizer-level instruction information for displaying hierarchy graph. Parent level is hlo level, child level is backend level. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``id`` - integer (int64) - - Yes - Unique identifier for corresponding assembly level instruction. * - ``raw_bir_id`` - integer (int64) - - Yes - Unique identifier for backend-level instruction. * - ``raw_penguin_id`` - integer (int64) - - Yes - Unique identifier for tensorizer-level instruction. * - ``raw_hlo_id`` - integer (int64) - - Yes - Unique identifier for hlo level instruction. * - ``raw_framework_id`` - integer (int64) - - Yes - Unique identifier for framework level instruction. * - ``penguin_name`` - string - - Yes - Tensorizer-level operation name. Example: ``float32<1 x 2048> $1345[i1_94_0, i1_94_1_0_0_0, i1_94_1_0_0_1]``. * - ``modular_flow_invocation_idx`` - integer (int64) - - No - Integer for identifying which call is the current instruction. 0 represent first time calling the function, 1 represent second time calling the function. Example: ``1``. * - ``start_ts`` - integer (int64) - nanoseconds - Yes - Minimum start time among all of the corresponding instruction timestamps. Unit is nanoseconds. * - ``end_ts`` - integer (int64) - nanoseconds - Yes - Maximum end time among all of the corresponding instruction timestamps. Unit is nanoseconds. * - ``time_intervals`` - string - - Yes - A list of timestamp intervals after merging all corresponding instruction timestamps from the lower hierarchy level. Each interval is represented as [start, end]. Example: ``[[1, 10], [20, 50]]``. * - ``element_count_per_interval`` - array of integer (int64) - - Yes - Array where each item corresponds to the total number of elements processed by a group of instructions from the lower hierarchy level. Each corresponds to the time interval in the same position in the time_intervals field. Example: ``[1024, 2048, 512]``. * - ``transpose_element_count_per_interval`` - array of integer (int64) - - Yes - Array where each item corresponds to the total number of transpose elements processed by a group of instructions from the lower hierarchy level. Each corresponds to the time interval in the same position in the time_intervals field. Example: ``[256, 512, 128]``. * - ``hbm_transfer_size_per_interval`` - array of integer (int64) - bytes - Yes - Array where each item corresponds to the total number of HBM bytes transferred by a group of instructions from the lower hierarchy level. Each corresponds to the time interval in the same position in the time_intervals field. Example: ``[16384, 32768, 8192]``. * - ``hbm_read_bytes`` - integer (int64) - bytes - Yes - Total number of bytes read from HBM by all grouped instructions in this tensorizer-level instruction. * - ``hbm_write_bytes`` - integer (int64) - bytes - Yes - Total number of bytes written to HBM by all grouped instructions in this tensorizer-level instruction. * - ``sbuf_read_bytes`` - integer (int64) - bytes - No - Total number of bytes read from SBUF by all grouped instructions in this tensorizer-level instruction. * - ``sbuf_write_bytes`` - integer (int64) - bytes - No - Total number of bytes written to SBUF by all grouped instructions in this tensorizer-level instruction. * - ``load_weight_bytes`` - integer (int64) - bytes - No - Total number of bytes loaded by LDWEIGHTS instructions in this tensorizer-level instruction. .. _profile-schema-psumusage: PsumUsage ~~~~~~~~~ PSUM memory bank utilization metrics during execution. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``timestamp`` - integer (int64) - nanoseconds - Yes - Timestamp for this usage measurement. * - ``avg_util_percent`` - number (float) - fraction (0-1) - No - Average fullness across all PSUM memory banks. Range: 0 to 1. * - ``bank_0_util_percent`` - number (float) - fraction (0-1) - No - Fullness percentage of PSUM memory bank 0. Range: 0 to 1. * - ``bank_1_util_percent`` - number (float) - fraction (0-1) - No - Fullness percentage of PSUM memory bank 1. Range: 0 to 1. * - ``bank_2_util_percent`` - number (float) - fraction (0-1) - No - Fullness percentage of PSUM memory bank 2. Range: 0 to 1. * - ``bank_3_util_percent`` - number (float) - fraction (0-1) - No - Fullness percentage of PSUM memory bank 3. Range: 0 to 1. * - ``bank_4_util_percent`` - number (float) - fraction (0-1) - No - Fullness percentage of PSUM memory bank 4. Range: 0 to 1. * - ``bank_5_util_percent`` - number (float) - fraction (0-1) - No - Fullness percentage of PSUM memory bank 5. Range: 0 to 1. * - ``bank_6_util_percent`` - number (float) - fraction (0-1) - No - Fullness percentage of PSUM memory bank 6. Range: 0 to 1. * - ``bank_7_util_percent`` - number (float) - fraction (0-1) - No - Fullness percentage of PSUM memory bank 7. Range: 0 to 1. * - ``pcore_idx`` - integer (int64) - - No - Index of the physical NeuronCore within a Logical NeuronCore (LNC). A Logical NeuronCore is a group of physical NeuronCores. For LNC2 this field can have a value of either 0 or 1. .. _profile-schema-sbufallocation: SbufAllocation ~~~~~~~~~~~~~~ State Buffer memory occupancy periods showing which specific address ranges are occupied over time. Uses memory lifetime analysis to track actual occupancy rather than individual operations. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``start_ts`` - integer (int64) - nanoseconds - Yes - Timestamp when the memory address range becomes occupied (allocated). * - ``end_ts`` - integer (int64) - nanoseconds - Yes - Timestamp when the memory address range becomes free (deallocated). * - ``start_address`` - integer (int64) - bytes - Yes - Starting memory address of the occupied region (inclusive - first byte of allocation). * - ``end_address`` - integer (int64) - bytes - Yes - Ending memory address of the occupied region (inclusive - last byte of allocation). * - ``partition`` - integer (int64) - - Yes - State Buffer partition number that this address range belongs to (0, 32, 64, or 96). * - ``op_source_type`` - string - - No - Source of memory operation. Distinguishes whether the memory allocation originated from a compute engine instruction or a DMA operation. Example: ``instruction``. * - ``opcode`` - string - - No - Instruction opcode name (only for instruction source type). Examples include MATMUL, TENSOR_TENSOR, etc. Example: ``MATMUL``. * - ``operands`` - string - - No - Instruction operands string (only for instruction source type). Contains source/destination addresses and other instruction parameters. Example: ``src=0x2120 dst=0x2312 rows=64 cols=64``. * - ``queue_name`` - string - - No - DMA queue name (only for DMA source type). Identifies which DMA queue performed the operation. Example: ``qGpSimdSpillReload0``. * - ``block_id`` - integer (int64) - - No - DMA block ID (only for DMA source type). Identifies which block within the DMA queue this operation corresponds to. * - ``tensor_name`` - string - - No - Name of the tensor or variable associated with this memory allocation (if available). Example: ``intermediate_result_0``. * - ``hlo_name`` - string - - No - High-Level Optimizer (HLO) operation name associated with this memory allocation (if available). Example: ``dot.1``. * - ``pcore_idx`` - integer (int64) - - No - Index of the physical NeuronCore within a Logical NeuronCore (LNC). A Logical NeuronCore is a group of physical NeuronCores. For LNC2 this field can have a value of either 0 or 1. .. _profile-schema-sbufusage: SbufUsage ~~~~~~~~~ State Buffer (SB) memory partition utilization metrics during execution. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``timestamp`` - integer (int64) - nanoseconds - Yes - Timestamp for this usage measurement. * - ``avg_util_percent`` - number (float) - fraction (0-1) - No - Average utilization across all State Buffer partitions. Range: 0 to 1. * - ``partition_0_util_percent`` - number (float) - fraction (0-1) - No - Fullness percentage of State Buffer partition 0. Range: 0 to 1. * - ``partition_32_util_percent`` - number (float) - fraction (0-1) - No - Fullness percentage of State Buffer partition 32. Range: 0 to 1. * - ``partition_64_util_percent`` - number (float) - fraction (0-1) - No - Fullness percentage of State Buffer partition 64. Range: 0 to 1. * - ``partition_96_util_percent`` - number (float) - fraction (0-1) - No - Fullness percentage of State Buffer partition 96. Range: 0 to 1. * - ``pcore_idx`` - integer (int64) - - No - Index of the physical NeuronCore within a Logical NeuronCore (LNC). A Logical NeuronCore is a group of physical NeuronCores. For LNC2 this field can have a value of either 0 or 1. .. _profile-schema-schemafields: SchemaFields ~~~~~~~~~~~~ The schema describes the data contained within the profile. Every field of every table is described in this table with details such as type, description, and more. See the "profile_schema_version" field of the Metadata table to determine the version of this schema. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``table_name`` - string - - Yes - The table this field belongs to. Example: ``Instruction``. * - ``table_description`` - string - - Yes - The description of the table this field belongs to. Example: ``Instructions executed on NeuronCore compute engines.``. * - ``field_name`` - string - - Yes - The name of this field. Example: ``engine``. * - ``field_description`` - string - - Yes - The description of this field. Example: ``Compute engine that executed this instruction.``. * - ``field_example`` - string - - No - One or more examples of what value this field may hold. Example: ``tensor``. * - ``field_type`` - string - - Yes - The data type of this field. Example: ``string``. * - ``field_required`` - boolean - - Yes - Indicates whether this field is required according to the OpenAPI schema definition. Example: ``true``. * - ``field_format`` - string - - No - The format provides additional detail about the data type of this field. The "string" and "boolean" types do not need a format, but other types such as "integer" can be more have a more specific format (such as "int64"). Example: ``int64``. * - ``field_unit`` - string - - No - The unit of this field. Example: ``nanoseconds``. * - ``field_minimum`` - number (float) - - No - The minimum value that can be held by this field. Example: ``0``. * - ``field_maximum`` - number (float) - - No - The maximum value that can be held by this field. Example: ``100``. * - ``field_enum_values`` - string - - No - Comma-separated list of allowed enum values for this field, if the field is an enum type. Example: ``instruction,input,output,data,software_dynamic,hardware_dynamic,unknown``. .. _profile-schema-semaphoreupdate: SemaphoreUpdate ~~~~~~~~~~~~~~~ Semaphore state changes used for synchronization between operations. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``timestamp`` - integer (int64) - nanoseconds - Yes - Timestamp when the semaphore update occurred. * - ``id`` - string - - Yes - Unique identifier of the semaphore being updated. * - ``value`` - integer (int32) - - Yes - New value assigned to the semaphore. * - ``update_source`` - string - - No - The source that generated this semaphore update notification (e.g. Tensor, Scalar, Vector, Sync, GpSimd, AXI). AXI indicates the update came from the AXI bus. Example: ``Tensor``. * - ``double_set_or_clear`` - boolean - - No - Indicates a double event set or clear condition. When true, the semaphore experienced a double-set (if value increased) or double-clear (if value decreased). Example: ``true``. * - ``update_mode`` - integer (int32) - - No - The raw semaphore update mode integer from the ISA TPB update mode enum. See update_operation_type and update_trigger_timing for decoded values. A value of 0 indicates no update (NONE). Example: ``0``. * - ``update_operation_type`` - string - - No - The decoded operation type from the update mode, describing what modification was applied to the semaphore or event. Possible values: none, sem_inc, sem_dec, sem_add_imm, sem_sub_imm, sem_wr_imm, evt_set, evt_clr, unknown. Example: ``sem_inc``. * - ``update_trigger_timing`` - string - - No - The decoded trigger timing from the update mode, indicating when the update fires relative to the instruction that requested it. Possible values: read (when the instruction begins reading its inputs) or complete (when the instruction finishes execution). Empty string when update_operation_type is none or unknown. Example: ``complete``. * - ``pcore_idx`` - integer (int64) - - No - Index of the physical NeuronCore within a Logical NeuronCore (LNC). A Logical NeuronCore is a group of physical NeuronCores. For LNC2 this field can have a value of either 0 or 1. .. _profile-schema-snapshotallocation: SnapshotAllocation ~~~~~~~~~~~~~~~~~~ Framework-level (PyTorch) tensor allocations captured by a memory snapshot (torch_neuronx record_memory_history). Each row is one allocation reconstructed by pairing an "alloc" trace event with its matching "free_completed" event by virtual address, in trace (ring-buffer) order. These represent the caching allocator's user-visible allocations in virtual address space. Recent runtime/torch changes put the snapshot timestamps into the system_clock domain, the SAME domain as the runtime ntrace.pb event timestamps, so alloc_ts/free_ts are the wall-clock x-axis for the Active Memory Timeline and align 1:1 with the system timeline. NOTE - this is distinct from the runtime-derived HbmAllocation table (physical HBM occupancy from the system trace) and from the runtime memory snapshot view. Use this table for the PyTorch "Active Memory Timeline". .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``alloc_ts`` - integer (int64) - microseconds - Yes - Timestamp (system_clock, ns-since-epoch / 1000) when the allocation occurred (alloc event). This is the timeline x-axis and aligns directly with ntrace.pb timestamp_ns / 1000. Pre-existing allocations (live before the trace window) are anchored to the window's first observed timestamp so their bands span the visible window. * - ``free_ts`` - integer (int64) - microseconds - Yes - Timestamp (system_clock, µs) when the allocation was freed (free_completed event). Set to the snapshot's last observed timestamp if never freed within the capture. * - ``alloc_index`` - integer (int64) - - No - Event-ordinal (0-based) of the alloc event in trace order. Retained as a stable secondary key: it encodes the ring-buffer replay order the stack topology depends on, breaks ties when two events share a timestamp, and serves as a fallback x-axis for older captures that predate the unified system_clock domain. Pre-existing allocations use 0 (window start). * - ``free_index`` - integer (int64) - - No - Event-ordinal of the free event. Equals the last event index when the allocation is never freed within the capture, so bands extend to the timeline end and no further. * - ``address`` - integer (int64) - bytes - Yes - Framework-visible virtual address of the allocation. * - ``size`` - integer (int64) - bytes - Yes - Requested allocation size in bytes. * - ``block_size`` - integer (int64) - bytes - No - Size of the underlying allocator block serving this allocation (may exceed size due to rounding). * - ``stream`` - integer (int64) - - No - Stream identifier associated with the allocation. * - ``device_idx`` - integer (int64) - - No - Device index the allocation belongs to (from the snapshot device_traces index). * - ``rank`` - integer (int64) - - Yes - Distributed rank of the process that produced this snapshot (parsed from memory_snapshot_rank.pickle). * - ``trace_id`` - integer (int64) - - No - Systrace correlation key linking this framework allocation to runtime system trace events. Correlation has three levels, all keyed on this id: (1) a snapshot alloc / segment_alloc trace_id=X matches the ntrace TENSOR_ALLOC event with tracking_id=X (already surfaced on HbmAllocation.trace_id); (2) ntrace TENSOR_READ / TENSOR_WRITE events carry a tensor_trace_id attribute that matches the same X, linking a framework tensor allocation to the runtime data transfers that touched it. 0 when not captured. * - ``clock_domain_offset_us`` - integer (int64) - - No - system_clock - steady_clock offset in microseconds, from the snapshot's clock_domain_offset_us field. alloc_ts/free_ts are already system_clock and need no conversion to align with ntrace.pb. This offset converts the framework profiler (profiler_trace.json Chrome-trace ts, which is steady_clock µs) into system_clock: profiler_ts + clock_domain_offset_us = system_clock_us. 0 for older captures that did not emit the field. Note - can be negative. Example: ``1710000000000000``. * - ``never_freed`` - boolean - - No - True when no matching free_completed event was observed within the capture (free_ts was set to the last observed timestamp). Example: ``false``. * - ``pre_existing`` - boolean - - No - True when this allocation was made before the trace ring-buffer window (either an unmatched free_completed, or an active_allocated segment block with no trace alloc). These form the timeline's baseline memory and are required for the peak to match PyTorch memory_viz. Example: ``false``. * - ``source_frames`` - string - - No - Newline-joined source call stack (innermost first) for the allocation. Example: ``rms_norm_neuron (rms_norm.py:18)``. .. _profile-schema-snapshotsegment: SnapshotSegment ~~~~~~~~~~~~~~~ Point-in-time view of the framework (PyTorch) caching allocator's internal state at the moment a memory snapshot was taken. Each row is one block within a segment (super-tensor allocated from NRT). Segments are cached and sub-divided into blocks; block state reveals which memory is actively used vs cached-but-idle (fragmentation). Segment-level fields are repeated on each of its block rows. NOTE - this is distinct from the runtime-derived HbmAllocation table and the SnapshotAllocation temporal table; it captures allocator caching state, not a timeline. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``segment_address`` - integer (int64) - bytes - Yes - Base virtual address of the segment (super-tensor). * - ``segment_total_size`` - integer (int64) - bytes - Yes - Total reserved size of the segment in bytes (cached from NRT, may exceed the sum of active blocks). * - ``segment_allocated_size`` - integer (int64) - bytes - No - Sum of currently-allocated block sizes within the segment. * - ``segment_active_size`` - integer (int64) - bytes - No - Sum of currently-active block sizes within the segment. * - ``is_small`` - boolean - - No - True if the segment belongs to the small-allocation pool (allocations below the small/large boundary), false for the large pool. Example: ``false``. * - ``stream`` - integer (int64) - - No - Stream identifier associated with the segment. * - ``device_idx`` - integer (int64) - - No - Device index the segment belongs to. * - ``rank`` - integer (int64) - - Yes - Distributed rank of the process that produced this snapshot (parsed from memory_snapshot_rank.pickle). * - ``block_address`` - integer (int64) - bytes - Yes - Virtual address of this block within the segment. * - ``block_size`` - integer (int64) - bytes - Yes - Size of this block in bytes (as carved by the allocator). * - ``block_requested_size`` - integer (int64) - bytes - No - Originally requested size for this block (may be smaller than block_size due to rounding). * - ``block_state`` - string - - Yes - Block state — active_allocated (in use), active_awaiting_free (freed by user, pending), or inactive (cached and free). Example: ``active_allocated``. * - ``source_frames`` - string - - No - Newline-joined source call stack (innermost first) of what last lived in this block, if available. Example: ``rms_norm_neuron (rms_norm.py:18)``. .. _profile-schema-stackframe: StackFrame ~~~~~~~~~~ Individual stack frames representing the call stack hierarchy for debugging and tracing execution back to source code. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``id`` - string - - Yes - Unique identifier for this stack frame (one-based index converted to string). * - ``file_location_id`` - integer (int32) - - Yes - Reference to the file location ID in StackFrameFileLocation table. * - ``parent_frame_id`` - string - - Yes - Reference to the parent stack frame ID, or 0 if this is the root frame. .. _profile-schema-stackframefilelocation: StackFrameFileLocation ~~~~~~~~~~~~~~~~~~~~~~ File location information combining file names, function names, and line numbers for debugging and tracing execution back to source code. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``id`` - string - - Yes - Unique identifier for this file location (one-based index converted to string). * - ``file_name_id`` - integer (int32) - - Yes - Reference to the file name ID in StackFrameFileName table. * - ``function_name_id`` - integer (int32) - - Yes - Reference to the function name ID in StackFrameFunctionName table. * - ``line_number`` - integer (int32) - - Yes - Line number within the source file. .. _profile-schema-stackframefilename: StackFrameFileName ~~~~~~~~~~~~~~~~~~ Source code file names referenced in the stack frames for debugging and tracing execution back to source code. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``id`` - string - - Yes - Unique identifier for this file name (one-based index converted to string). * - ``name`` - string - - Yes - Source code file name or path. Example: ``model.py``. .. _profile-schema-stackframefunctionname: StackFrameFunctionName ~~~~~~~~~~~~~~~~~~~~~~ Function names referenced in the stack frames for debugging and tracing execution back to source code. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``id`` - string - - Yes - Unique identifier for this function name (one-based index converted to string). * - ``name`` - string - - Yes - Function or method name. .. _profile-schema-summary: Summary ~~~~~~~ Profile performance metrics which summarize execution. For profile metadata and system information see the Metadata table. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``total_time`` - number (float) - seconds - No - Total duration of on device execution time for the run, as reported by the runtime from explicit TPB start/stop notifications stored in the NTFF. This is the hardware-measured wall-clock time from execution start to execution end, including any idle gaps between instructions. It does not include host-side overhead such as host<->device data movement or host runtime/framework overhead. When explicit notifications are unavailable, falls back to the span of all parsed trace notifications. Units are in seconds. * - ``neuroncore_cycle_count`` - integer (int64) - cycles - No - Number of NeuronCore cycles over the duration of the profile i.e. the duration in seconds multiplied by the NeuronCore base clock rate The NeuronCore base clock rate for Inferentia1 is 1.0 GHz, Trainium1 is 1.4 GHz, and Trainium2 is 1.2 GHz. * - ``model_flops`` - integer (int64) - flops - No - Model FLOPs is the FLOP count of useful compute on the Tensor Engine i.e. the matrix multiplications from the HLO model definition, and excludes other instructions for data movement (i.e. transposes and partition broadcasts) inserted by the compiler to resolve memory layout conflicts. Note, each floating point multiply-add is counted as two FLOPs. Calculated as 2 \* MAC_count from the HLO stats. * - ``hardware_flops`` - integer (int64) - flops - No - Hardware FLOPs is the FLOP count calculated from all Tensor Engine instructions that Neuron Compiler emits for execution. This metric is typically higher than model_flops because it includes matmul instructions for data movement (i.e. transposes and partition broadcasts) inserted by the compiler to resolve memory layout conflicts. Note, each floating point multiply-add is counted as two FLOPs. Calculated as the sum of (2 \* tensor_matrix_active_rows \* tensor_matrix_active_cols \* elements) for each matrix multiplication. * - ``transpose_flops`` - integer (int64) - flops - No - 2x the number of MATMUL operations from transposes. A MATMUL operation is counted as a transpose if the instruction type in its debug info is "TRANSPOSE". This is a subset of hardware_flops. * - ``adjusted_hardware_flops`` - integer (int64) - flops - No - BF16-equivalent hardware FLOPs. Scales FLOPs to the BF16 throughput equivalent so the result can be used with BF16 peak throughput as the denominator for PE utilization calculations. BF16, FP16, and TF32 use a 1.0x multiplier. FP32 uses 4.0x. On Trn2/Trn3, FP8@FP8 and UINT8@UINT8 matmuls use 0.5x; FP8 or UINT8 mixed with a higher-precision operand uses 1.0x (executes at BF16 throughput). FP8 uses 1.0x on Trn1/Inf2. MXFP8 and MXFP4 use 0.25x on Trn3. * - ``adjusted_transpose_flops`` - integer (int64) - flops - No - BF16-equivalent transpose FLOPs. Scales transpose FLOPs to the BF16 throughput equivalent so the result can be used with BF16 peak throughput as the denominator for PE utilization calculations. Uses the same scaling multipliers as adjusted_hardware_flops. * - ``mm_arithmetic_intensity`` - number (float) - - No - The ratio of regular MATMUL flops to total DRAM transfer size. If peak_flops_bandwidth_ratio is greater than this value, the workload is memory bound. If it is less than this value, the workload is compute bound. It is calculated as (hardware_flops - transpose_flops) / (hbm_write_bytes + hbm_read_bytes). * - ``peak_flops_bandwidth_ratio`` - number (float) - - No - The ratio of theoretical max Tensor Engine FLOPS to peak DRAM bandwidth. If mm_arithmetic_intensity is less than this value, the workload is memory bound. If it is greater than this value, the workload is compute bound. * - ``trace_count`` - integer (int64) - - No - Total number of trace events in the NTFF including instructions, events/semaphores, and DMAs. * - ``event_count`` - integer (int64) - - No - Total number of event notifications, such as for semaphore updates. This count is less than trace_count because it does not count instructions. * - ``total_row_count`` - integer (int64) - - No - Total number of data points written to the database. * - ``dma_active_cycles`` - integer (int64) - cycles - No - dma_active_time multiplied by TPB clock speed. * - ``dma_active_time`` - number (float) - seconds - No - Duration of time when at least one DMA packet is transfering. * - ``dma_active_time_percent`` - number (float) - fraction (0-1) - No - Percentage of time when at least one DMA packet is transfering. Range: 0 to 1. * - ``host_device_dma_packet_count`` - integer (int64) - - No - Number of host-device DMA packets during the profiled execution window. The runtime and driver call this H2D DMA. * - ``host_device_dma_transfer_bytes`` - integer (int64) - bytes - No - Total bytes transferred by host-device DMA packets in either direction during the profiled execution window. * - ``host_device_dma_active_time`` - number (float) - seconds - No - Duration when at least one host-device DMA engine was transferring data during the profiled execution window. * - ``host_device_dma_bandwidth`` - number (float) - bytes_per_second - No - Host-device DMA transfer bytes divided by the duration of the profiled execution window. * - ``static_dma_active_time`` - number (float) - seconds - No - Duration of time when at least one statically generated DMA packet is transfering. * - ``dynamic_dma_active_time_percent`` - number (float) - fraction (0-1) - No - Percentage of time when at least one dynamically generated DMA packet is transfering. Range: 0 to 1. * - ``static_dma_active_time_percent`` - number (float) - fraction (0-1) - No - Percentage of time when at least one statically generated DMA packet is transfering. Range: 0 to 1. * - ``dma_transfer_average_bytes`` - number (float) - bytes - No - Total transfer size of all DMA packets divided by the total number of DMA transfers. A DMA transfer is a high level compiler construct representing the movement of one or more tensors, while DMA packets are standard DMA terminology and are the smallest unit of DMA data movement displayed in the Neuron Explorer timeline. A single DMA transfer is fulfilled by multiple DMA packets. The Neuron Explorer timeline displays DMA Packets when "Show expanded DMA" is ON, otherwise it aggregates them into DMA transfers. * - ``dma_packet_time`` - number (float) - seconds - No - Total duration of all DMA engine packets. Overlapping DMA packets across all DMA engines are added together. * - ``dma_transfer_time`` - number (float) - seconds - No - Total duration of all DMA transfers. A DMA transfer is a high level compiler construct representing the movement of one or more tensors, while DMA packets are standard DMA terminology and are the smallest unit of DMA data movement displayed in the Neuron Explorer timeline. A single DMA transfer is fulfilled by multiple DMA packets. The Neuron Explorer timeline displays DMA Packets when "Show expanded DMA" is ON, otherwise it aggregates them into DMA transfers. * - ``dma_transfer_count`` - integer (int64) - - No - Number of DMA transfers. A DMA transfer is a high level compiler construct representing the movement of one or more tensors, while DMA packets are standard DMA terminology and are the smallest unit of DMA data movement displayed in the Neuron Explorer timeline. A single DMA transfer is fulfilled by multiple DMA packets. The Neuron Explorer timeline displays DMA Packets when "Show expanded DMA" is ON, otherwise it aggregates them into DMA transfers. * - ``dma_transfer_total_bytes`` - integer (int64) - bytes - No - Total transfer size of all DMA transfers. A DMA transfer is a high level compiler construct representing the movement of one or more tensors. * - ``dma_queue_count`` - integer (int64) - - No - Total number of DMA queues the profile reports, including all input, output, and weight DMA queues for different DMA engines, and the collectives queues. * - ``inputs_and_weights_size_bytes`` - integer (int64) - bytes - No - Size of input tensors and weight tensors. * - ``inputs_outputs_weights_size_bytes`` - integer (int64) - bytes - No - Size of input tensors, output tensors, and weight tensors. Used to support perf bounds, specifically MemoryBoundIdealNoReloads * - ``weight_size_bytes`` - integer (int64) - bytes - No - Total size of all weights defined in the NEFF. * - ``static_dma_packet_count`` - integer (int64) - - No - Number of DMA packets that are statically generated. * - ``static_dma_size`` - integer (int64) - bytes - No - Total transfer size of DMA packets that are statically generated. * - ``dynamic_dma_packet_percent`` - number (float) - fraction (0-1) - No - Percentage of DMA packets that are dynamically generated during execution (generally higher means better performance). Range: 0 to 1. * - ``static_dma_packet_percent`` - number (float) - fraction (0-1) - No - Percentage of DMA packets that are statically generated. Range: 0 to 1. * - ``dynamic_dma_size_percent`` - number (float) - fraction (0-1) - No - Percentage of total transfer size for DMA packets that are dynamically generated during execution. Range: 0 to 1. * - ``static_dma_size_percent`` - number (float) - fraction (0-1) - No - Percentage of total transfer size for DMA packets that are statically generated. Range: 0 to 1. * - ``mbu_estimated_percent`` - number (float) - fraction (0-1) - No - Memory Bandwidth Utilization (MBU) is the achieved HBM bandwidth utilization. It is estimated as a ratio of achieved_hbm_bandwidth over hbm_max_bandwidth, where achieved_hbm_bandwidth is total bytes read/written from/to HBM divided by total_time and hbm_max_bandwidth. Calculated as (hbm_write_bytes + hbm_read_bytes) / total_time / hbm_max_bandwidth. The hbm_max_bandwidth is the maximum HBM bandwidth of a HBM stack. For LNC1 and LNC2, the hbm_max_bandwidth is the same. Range: 0 to 1. * - ``mbu_min_read_util_percent`` - number (float) - fraction - No - Estimated minimum possible Memory Bandwidth Utilization (MBU) utilization assuming inputs and weights are read from HBM a single time (no repeat loads). Calculated as inputs_and_weights_size_bytes / total_time / hbm_max_bandwidth. This value can exceed 1.0, because inputs_and_weights_size_bytes is the declared size of the NEFF's inputs and weights, which can be more data than could possibly be read in the elapsed time of the profile. Use mbu_estimated_percent for utilization based only on bytes observed on the HBM bus. mbu_min_read_util_percent may not be valid if not all inputs and weights in the NEFF are loaded, so the value can be misleading below 1.0 as well. * - ``hbm_write_bytes`` - integer (int64) - bytes - No - Total bytes of data written to HBM using the DMA engines. This does not include data movement from host DRAM to HBM because this happens before the execution profiling session begins. * - ``hbm_read_bytes`` - integer (int64) - bytes - No - Total bytes of data read from HBM using the DMA engines. This does not include data movement from HBM to host DRAM because this happens after the execution profiling session. * - ``sbuf_write_bytes`` - integer (int64) - bytes - No - Total size of all writes to the State Buffer. This includes DMAs writing to and instructions with output to the State Buffer. * - ``sbuf_read_bytes`` - integer (int64) - bytes - No - Total size of all reads from the State Buffer. This includes DMAs reading from and instructions with input from the State Buffer. * - ``spill_save_bytes`` - integer (int64) - bytes - No - Total bytes of spilled data that was saved to HBM. Spilled data is the intermediate tensors computed by the engines that cannot fit in the SBUF during execution and must be spilled into HBM. * - ``spill_reload_bytes`` - integer (int64) - bytes - No - Total bytes of spilled data that was reloaded back to SBUF. Spilled data is the intermediate tensors computed by the engines that cannot fit in the SBUF during execution and must be spilled into HBM. If a spilled tensor is reloaded multiple times into SBUF, this metric will include the spilled tensor size multiplied by the reload count. * - ``psum_write_bytes`` - integer (int64) - bytes - No - Total bytes of data that are written to PSUM by compute engine instructions. * - ``psum_read_bytes`` - integer (int64) - bytes - No - Total bytes of data that are read from PSUM by compute engine instructions. * - ``psum_read_sbuf_write_count`` - integer (int64) - - No - Number of compute engine instructions that moved data from PSUM to SBUF. * - ``psum_read_sbuf_write_bytes`` - integer (int64) - bytes - No - Total bytes of data moved from PSUM to SBUF by compute engine instructions. * - ``input_queue_bytes`` - integer (int64) - bytes - No - Total transfer size of DMA blocks in "input" queues. The "input" queues are primarily for loading input data from HBM to the State Buffer. * - ``output_queue_bytes`` - integer (int64) - bytes - No - Total transfer size of DMA blocks in "output" queues. The "output" queues are primarily for writing output data from the State Buffer to HBM. * - ``weight_queue_bytes`` - integer (int64) - bytes - No - Total transfer size of DMA blocks in "data" queues. The "data" queues are primarily for loading weights from HBM to the State Buffer. This size may be larger than weight_size_bytes if weights are loaded multiple times or spilled and reloaded. * - ``software_dynamic_dma_active_time`` - number (float) - seconds - No - Duration of time when at least one software generated DMA packet is transfering. * - ``software_dynamic_dma_active_time_percent`` - number (float) - fraction (0-1) - No - Percentage of time when at least one software generated DMA packet is transfering. Range: 0 to 1. * - ``software_dynamic_dma_packet_count`` - integer (int64) - - No - Number of DMA packets that are dynamically generated during execution via software backend. * - ``software_dynamic_dma_size`` - integer (int64) - bytes - No - Total transfer size of DMA packets that are dynamically generated during execution via software backend. * - ``software_dynamic_dma_packet_percent`` - number (float) - fraction (0-1) - No - Percentage of DMA packets that are dynamically generated during execution via software backend. Range: 0 to 1. * - ``software_dynamic_dma_size_percent`` - number (float) - fraction (0-1) - No - Percentage of total transfer size for DMA packets that are dynamically generated via software backend. Range: 0 to 1. * - ``hardware_dynamic_dma_active_time`` - number (float) - seconds - No - Duration of time when at least one hardware generated DMA packet is transfering. * - ``hardware_dynamic_dma_active_time_percent`` - number (float) - fraction (0-1) - No - Percentage of time when at least one hardware generated DMA packet is transfering. Range: 0 to 1. * - ``hardware_dynamic_dma_packet_count`` - integer (int64) - - No - Number of DMA packets that are dynamically generated during execution via dedicated hardware. * - ``hardware_dynamic_dma_size`` - integer (int64) - bytes - No - Total transfer size of DMA packets that are dynamically generated during execution via dedicated hardware. * - ``hardware_dynamic_dma_packet_percent`` - number (float) - fraction (0-1) - No - Percentage of DMA packets that are dynamically generated during execution via dedicated hardware. Range: 0 to 1. * - ``hardware_dynamic_dma_size_percent`` - number (float) - fraction (0-1) - No - Percentage of total transfer size for DMA packets that are dynamically generated via dedicated hardware. Range: 0 to 1. * - ``tensor_engine_instruction_time`` - number (float) - seconds - No - Total duration of all Tensor engine instructions. Overlapping instructions on the engine are added together. Units are in seconds. * - ``tensor_engine_active_time`` - number (float) - seconds - No - Duration of time when Tensor engine is processing at least one instruction (excluding semaphore waits). Units are in seconds. * - ``tensor_engine_active_time_percent`` - number (float) - fraction (0-1) - No - Percentage of time when Tensor engine is processing at least one instruction (excluding semaphore waits). Range: 0 to 1. * - ``tensor_engine_instruction_count`` - integer (int64) - - No - Number of Tensor Engine instructions. * - ``matmul_instruction_count`` - integer (int64) - - No - Total number of MATMUL instructions. * - ``mfu_estimated_percent`` - number (float) - fraction (0-1) - No - MFU is Model FLOPs Utilization. This reflects the Tensor Engine utilization for useful compute i.e. the matrix multiplications from the HLO model definition, and excludes other instructions for data movement (i.e. transposes and partition broadcasts) inserted by the compiler to resolve memory layout conflicts. Note, each floating point multiply-add is counted as two FLOPs. Calculated as 2 \* MAC_count / (tensor_engine_max_ops_per_sec \* total_time) where tensor_engine_max_ops_per_sec is 2 times the number of Tensor Engine elements times the clock speed. Range: 0 to 1. * - ``mfu_max_achievable_estimated_percent`` - number (float) - fraction (0-1) - No - The best Tensor Engine Utilization you can achieve for this workload. Calculated as (model_flops / (hbm_write_bytes + hbm_read_bytes)) / (max_ops_per_sec / max_hbm_bandwidth). If mfu_estimated_percent is much smaller than this value, then this likely indicates a major data movement inefficiency introduced by compilation or kernel implementation for NKI. If mfu_estimated_percent is close to this value and the mfu_estimated_percent value is also very high, then there's no point reducing HBM traffic. But if the values are close and mfu_estimated_percent is lower than you want, consider improving the arithmetic intensity of your workload by performing more computation on every data access. Range: 0 to 1. * - ``mfu_hlo_estimated_percent`` - number (float) - fraction (0-1) - No - MFU (HLO) is Model FLOPs Utilization based on HLO stats. This reflects the Tensor Engine utilization for useful compute i.e. the matrix multiplications from the HLO model definition, and excludes other instructions for data movement (i.e. transposes and partition broadcasts) inserted by the compiler to resolve memory layout conflicts. Note, each floating point multiply-add is counted as two FLOPs. Calculated as model_flops / (tensor_engine_max_ops_per_sec \* total_time) where model_flops is 2 \* MAC_count from HLO stats and tensor_engine_max_ops_per_sec is 2 times the number of Tensor Engine elements times the clock speed. Note this will be 0 if HLO stats are missing, which is the case for NKI kernels. Range: 0 to 1. * - ``mfu_hlo_max_achievable_estimated_percent`` - number (float) - fraction (0-1) - No - The best Tensor Engine Utilization you can achieve for this workload based on HLO model FLOPs. Calculated as (model_flops / (hbm_write_bytes + hbm_read_bytes)) / (max_ops_per_sec / max_hbm_bandwidth). If mfu_hlo_estimated_percent is much smaller than this value, then this likely indicates a major data movement inefficiency introduced by compilation or kernel implementation for NKI. If mfu_hlo_estimated_percent is close to this value and the mfu_hlo_estimated_percent value is also very high, then there's no point reducing HBM traffic. But if the values are close and mfu_hlo_estimated_percent is lower than you want, consider improving the arithmetic intensity of your workload by performing more computation on every data access. Note this will be 0 if HLO stats are missing, which is the case for NKI kernels. Range: 0 to 1. * - ``mfu_inst_estimated_percent`` - number (float) - fraction (0-1) - No - MFU (Instruction) is Model FLOPs Utilization based on instruction trace. This reflects the Tensor Engine utilization for useful compute calculated from executed instructions, excluding transposes. Unlike mfu_hlo_estimated_percent which uses HLO model stats, this metric uses (hardware_flops - transpose_flops) from the actual instruction trace. Note, each floating point multiply-add is counted as two FLOPs. Calculated as (hardware_flops - transpose_flops) / (tensor_engine_max_ops_per_sec \* total_time) where tensor_engine_max_ops_per_sec is 2 times the number of Tensor Engine elements times the clock speed. Range: 0 to 1. * - ``mfu_inst_max_achievable_estimated_percent`` - number (float) - fraction (0-1) - No - The best Tensor Engine Utilization you can achieve for this workload based on instruction trace FLOPs (excluding transposes). Calculated as ((hardware_flops - transpose_flops) / (hbm_write_bytes + hbm_read_bytes)) / (max_ops_per_sec / max_hbm_bandwidth). If mfu_inst_estimated_percent is much smaller than this value, then this likely indicates a major data movement inefficiency. If mfu_inst_estimated_percent is close to this value and the mfu_inst_estimated_percent value is also very high, then there's no point reducing HBM traffic. But if the values are close and mfu_inst_estimated_percent is lower than you want, consider improving the arithmetic intensity of your workload by performing more computation on every data access. Range: 0 to 1. * - ``hfu_estimated_percent`` - number (float) - fraction (0-1) - No - HFU is Hardware FLOPs Utilization. This reflects the Tensor Engine utilization calculated from all Tensor Engine instructions that Neuron Compiler emits for execution. This metric is typically higher than mfu_estimated_percent because it includes matmul instructions for data movement (i.e. transposes and partition broadcasts) inserted by the compiler to resolve memory layout conflicts. Note, each floating point multiply-add is counted as two FLOPs. Calculated as hardware_flops / (tensor_engine_max_ops_per_sec \* total_time) where tensor_engine_max_ops_per_sec is 2 times the number of Tensor Engine elements times the clock speed. Range: 0 to 1. * - ``vector_engine_instruction_time`` - number (float) - seconds - No - Total duration of all Vector engine instructions. Overlapping instructions on the engine are added together. Units are in seconds. * - ``vector_engine_active_time`` - number (float) - seconds - No - Duration of time when Vector engine is processing at least one instruction (excluding semaphore waits). Units are in seconds. * - ``vector_engine_active_time_percent`` - number (float) - fraction (0-1) - No - Percentage of time when Vector engine is processing at least one instruction (excluding semaphore waits). Range: 0 to 1. * - ``vector_engine_instruction_count`` - integer (int64) - - No - Number of Vector Engine instructions. * - ``scalar_engine_instruction_time`` - number (float) - seconds - No - Total duration of all Scalar engine instructions. Overlapping instructions on the engine are added together. Units are in seconds. * - ``scalar_engine_active_time`` - number (float) - seconds - No - Duration of time when Scalar engine is processing at least one instruction (excluding semaphore waits). Units are in seconds. * - ``scalar_engine_active_time_percent`` - number (float) - fraction (0-1) - No - Percentage of time when Scalar engine is processing at least one instruction (excluding semaphore waits). Range: 0 to 1. * - ``scalar_engine_instruction_count`` - integer (int64) - - No - Number of Scalar Engine instructions. * - ``activate_instruction_time`` - number (float) - seconds - No - Total duration of all ACTIVATE/ACTIVATE_QUANTIZE instructions. Overlapping instructions are added up. Units are in seconds. * - ``activate_instruction_count`` - integer (int64) - - No - Number of ACTIVATE/ACTIVATE_QUANTIZE instructions. * - ``sync_engine_instruction_time`` - number (float) - seconds - No - Total duration of all Sync engine instructions. Overlapping instructions on the engine are added together. Units are in seconds. * - ``sync_engine_active_time`` - number (float) - seconds - No - Duration of time when Sync engine is processing at least one instruction (excluding semaphore waits). Units are in seconds. * - ``sync_engine_active_time_percent`` - number (float) - fraction (0-1) - No - Percentage of time when Sync engine is processing at least one instruction (excluding semaphore waits). Range: 0 to 1. * - ``sync_engine_instruction_count`` - integer (int64) - - No - Number of Sync Engine instructions. * - ``gpsimd_engine_instruction_time`` - number (float) - seconds - No - Total duration of all GpSimd engine instructions. Overlapping instructions on the engine are added together. Units are in seconds. * - ``gpsimd_engine_active_time`` - number (float) - seconds - No - Duration of time when GpSimd engine is processing at least one instruction (excluding semaphore waits). Units are in seconds. * - ``gpsimd_engine_active_time_percent`` - number (float) - fraction (0-1) - No - Percentage of time when GpSimd engine is processing at least one instruction (excluding semaphore waits). Range: 0 to 1. * - ``gpsimd_engine_instruction_count`` - integer (int64) - - No - Number of GpSimd Engine instructions. * - ``cc_op_time`` - number (float) - seconds - No - Total duration of all Collective Communication (CC) operations. Overlapping operations are added together. Units are in seconds. * - ``cc_op_active_time`` - number (float) - seconds - No - Duration of time when at least one Collective Communication (CC) operation is active. Units are in seconds. * - ``cc_op_active_time_percent`` - number (float) - fraction (0-1) - No - Percentage of time when at least one Collective Communication (CC) operation is active. Range: 0 to 1. * - ``cc_op_count`` - integer (int64) - - No - Number of Collective Communication (CC) operations. * - ``cc_cores_instruction_time`` - number (float) - seconds - No - Total duration of instructions on all Collective Communication (CC) cores. Overlapping instructions on the cores are added together. * - ``cc_cores_instruction_active_time`` - number (float) - seconds - No - Duration of time when at least one Collective Communication (CC) core is processing an instruction (excluding semaphore waits). * - ``cc_cores_instruction_active_time_percent`` - number (float) - fraction (0-1) - No - Percentage of time when at least one Collective Communication (CC) core is processing an instruction (excluding semaphore waits). Range: 0 to 1. * - ``cc_cores_instruction_count`` - integer (int64) - - No - Number of instructions on all Collective Communication (CC) cores. * - ``total_exec_time`` - number (float) - seconds - No - Duration from the earliest non-setup instruction or DMA packet to the latest, excluding setup/teardown opcodes (NOP, SET_ORDERING_MODE, EVENT_SEMAPHORE, EVENT_SEMAPHORE_RANGE_CLEAR, NOTIFY, COMPARE_BRANCH, DRAIN, WRITE, TENSOR_LOAD, MODIFY_POOL_CONFIG). This measures the wall-clock span of active inference work on the device. Units are in seconds. * - ``total_active_time`` - number (float) - seconds - No - Total duration of time when at least one engine, DMA, or CC operation is active. Computed as the union of all active intervals. Units are in seconds. * - ``total_active_time_percent`` - number (float) - fraction (0-1) - No - Percentage of total execution time when at least one engine, DMA, or CC operation is active. Range: 0 to 1. .. _profile-schema-systemprofileevents: SystemProfileEvents ~~~~~~~~~~~~~~~~~~~ This table contains system profile events, which are generated from the Neuron Runtime, the framework code, or the Python application. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``id`` - integer (int64) - - Yes - Unique identifier for this event (framework event ids are uint64 with bit 62 set, so they exceed JavaScript's safe-integer range). Served as native int64 over Arrow and read in the UI as BigInt, preserving full precision — no VARCHAR cast. generate-schema.py emits this field as bigint in the TypeScript schema. * - ``name`` - string - - Yes - Name of the event. Example: ``nrt_execute``. * - ``start_ts`` - integer (int64) - nanoseconds - Yes - Timestamp when event started. * - ``end_ts`` - integer (int64) - nanoseconds - Yes - Timestamp when event finished. * - ``duration_ns`` - integer (int64) - nanoseconds - Yes - Duration of this event. * - ``trace_event_source`` - string - - Yes - Source of the event. The value "neuron_rt" means Neuron Runtime running on CPU produced this event. The value "neuron_hw" means a Neuron Device produced this event. The value "framework" means the Python application or framework code produced this event. Example: ``["neuron_rt", "neuron_hw", "framework"]``. * - ``instance_id`` - string - - Yes - Unique identifier of the server instance that produced this event. Example: ``i-0b1ea78ca2865fd32``. * - ``process_id`` - integer (int64) - - Yes - The process id on that produced this event. Not necessarily unique when multiple instances or containers are in use. * - ``thread_id`` - integer (int64) - - Yes - The thread id on that produced this event. Not necessarily unique when multiple instances or containers are in use. * - ``nc_idx`` - integer (int64) - - No - Index of the Logical NeuronCore (LNC) that produced this event. Present on runtime and hardware events that are associated with a specific NeuronCore. * - ``pcore_idx`` - integer (int64) - - No - Index of the physical NeuronCore within a NeuronDevice. * - ``worker_gid`` - integer (int64) - - No - Global worker identifier that uniquely identifies a NeuronCore worker across all instances in a workload. Unlike nc_idx (which is process-local), worker_gid is unique across all processes and instances. * - ``flow_id`` - array of string - - No - Optional list of flow IDs associated with this event. Only present on events that participate in a flow (e.g. linking a framework dispatch event to the corresponding runtime execution event). Values are string-encoded integers to preserve precision in JSON and JavaScript. Example: ``["100", "200"]``. * - ``framework_stream_id`` - string - - No - Framework stream identifier. For framework events with category privateuse1_driver, this is the thread ID of the driver event. For neuron_rt events, this is propagated from the linked driver event via nrt_sequence_id. Used to sub-group events by stream in the timeline viewer. Example: ``0``. * - ``extra_attributes_json`` - string - - No - JSON string containing any additional attributes that were present in the event but don't have dedicated column in this table yet. This allows capturing unrecognized fields without data loss. Do not depend on this field because key/values here will move to dedicated columns. Columns are better because they are documented, typed correctly, higher performance, and more searchable. Example: ``{"some_field": "value", "another_field": 123}``. .. _profile-schema-systemprofilehbmusage: SystemProfileHbmUsage ~~~~~~~~~~~~~~~~~~~~~ Neuron Device HBM memory utilization metrics in bytes by usage type. Each row represents HBM utilization at a specific timestamp with all usage types aggregated into a single row for stacked area chart visualization. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``timestamp`` - integer (int64) - nanoseconds - Yes - Timestamp for this HBM memory utilization measurement. * - ``total_bytes`` - integer (int64) - bytes - Yes - Total HBM bytes used across all usage types. * - ``generic_bytes`` - integer (int64) - bytes - Yes - HBM bytes used for Generic allocations. * - ``tensor_instructions_bytes`` - integer (int64) - bytes - Yes - HBM bytes used for Tensor Instructions. * - ``scalar_instructions_bytes`` - integer (int64) - bytes - Yes - HBM bytes used for Scalar Instructions. * - ``gpsimd_instructions_bytes`` - integer (int64) - bytes - Yes - HBM bytes used for GpSimd Instructions. * - ``io_bytes`` - integer (int64) - bytes - Yes - HBM bytes used for IO. * - ``dram_spill_bytes`` - integer (int64) - bytes - Yes - HBM bytes used for DRAM Spill. * - ``weights_bytes`` - integer (int64) - bytes - Yes - HBM bytes used for Weights. * - ``profiler_buffers_bytes`` - integer (int64) - bytes - Yes - HBM bytes used for Profiler Buffers. * - ``scalar_table_bytes`` - integer (int64) - bytes - Yes - HBM bytes used for Scalar Table. * - ``vector_instructions_bytes`` - integer (int64) - bytes - Yes - HBM bytes used for Vector Instructions. * - ``sync_instructions_bytes`` - integer (int64) - bytes - Yes - HBM bytes used for Sync Instructions. * - ``shared_scratchpad_bytes`` - integer (int64) - bytes - Yes - HBM bytes used for Shared Scratchpad. * - ``tensors_bytes`` - integer (int64) - bytes - Yes - HBM bytes used for Tensors. * - ``ucode_lib_bytes`` - integer (int64) - bytes - Yes - HBM bytes used for Ucode lib. * - ``gpsimd_stdio_bytes`` - integer (int64) - bytes - Yes - HBM bytes used for GpSimd STDIO. * - ``collectives_bytes`` - integer (int64) - bytes - Yes - HBM bytes used for Collectives. * - ``scratchpad_bytes`` - integer (int64) - bytes - Yes - HBM bytes used for Scratchpad. * - ``dma_rings_runtime_bytes`` - integer (int64) - bytes - Yes - HBM bytes used for DMA Rings Runtime. * - ``dma_rings_spill_bytes`` - integer (int64) - bytes - Yes - HBM bytes used for DMA Rings Spill. * - ``dma_rings_io_bytes`` - integer (int64) - bytes - Yes - HBM bytes used for DMA Rings IO. * - ``dma_rings_collectives_bytes`` - integer (int64) - bytes - Yes - HBM bytes used for DMA Rings Collectives. * - ``unknown_bytes`` - integer (int64) - bytes - Yes - HBM bytes used for unrecognized usage types. * - ``device_idx`` - integer (int64) - - Yes - Neuron device index relative to the instance * - ``hbm_idx`` - integer (int64) - - Yes - The HBM index relative to the device. Depending on the architecture and logical NeuronCore configuration, multiple NeuronCores may share the same HBM. * - ``instance_id`` - string - - Yes - Identifier for the instance on which this HBM memory utilization was measured. Example: ``i-1234567890abcdef0``. .. _profile-schema-systemprofilemetadata: SystemProfileMetadata ~~~~~~~~~~~~~~~~~~~~~ Metadata for a system profile. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``first_ts`` - string - - Yes - First timestamp in the system profile trace, in human-readable format. Example: ``2025-12-18 15:35:56.592459482 +0000 UTC``. * - ``last_ts`` - string - - Yes - Last timestamp in the system profile trace, in human-readable format. Example: ``2025-12-18 15:35:57.340255583 +0000 UTC``. * - ``first_ts_ns`` - string - - Yes - First timestamp of the system profile trace, in nanoseconds, relative to the zeroed timeline origin. Always 0 for profiles processed by the current backend; older profiles carry an absolute Unix-nanosecond value that the UI rebases. Example: ``0``. * - ``last_ts_ns`` - string - - Yes - Last timestamp of the system profile trace, in nanoseconds, relative to the zeroed timeline origin (i.e. the trace duration). Older profiles carry an absolute Unix-nanosecond value. Example: ``747796101``. * - ``ntff_version`` - integer (int64) - - Yes - Version of the NTFF (Neuron Trace File Format) used for this profile. Example: ``5``. * - ``hbm_capacity_bytes`` - integer (int64) - bytes - Yes - Total HBM capacity in bytes for a single NeuronCore. * - ``instance_type`` - string - - Yes - The instance type on which this profile was captured. Example: ``trn1.32xlarge``. .. _profile-schema-tensorinfo: TensorInfo ~~~~~~~~~~ Information about tensors used in the NEFF. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``variable_name`` - string - - No - The name of the tensor. Example: ``input0``. * - ``kernel_param_name`` - string - - No - The NKI kernel parameter name for this tensor, if available. Provides a human-readable name that describes the tensor's role within the kernel (e.g. "expert_gate_up_weights" instead of the generic "input17"). Example: ``expert_gate_up_weights``. * - ``type`` - string - - No - What the tensor is used for. For example input tensor, output tensor, or weight tensor. Example: ``["IN", "OUT", "WEIGHT"]``. * - ``format`` - string - - No - Layout of the tensor in memory. For example, "NHWC" indicates a specific arrangement of dimensions. Letters you may see include N (batch size), H (height), W (width), C (channel). Example: ``NHWC``. * - ``shape`` - string - - No - The multi-dimensional shape of the tensor. Example: ``["128 128 1 1"]``. * - ``size`` - integer (int64) - bytes - No - The total size of the tensor in bytes. * - ``node`` - integer (int64) - - No - Neff node. Example: ``50``. * - ``pcore_idx`` - integer (int64) - - No - Index of the physical NeuronCore within a Logical NeuronCore (LNC). A Logical NeuronCore is a group of physical NeuronCores. For LNC2 this field can have a value of either 0 or 1. * - ``load_to_sbuf_avg_size_bytes`` - integer (int64) - - No - The average size in bytes of each DMA transfer when loading this tensor into the State Buffer. * - ``load_to_sbuf_total_size_bytes`` - integer (int64) - - No - The total size in bytes of all DMA transfers when loading this tensor into the State Buffer. * - ``load_to_sbuf_dma_count`` - integer (int64) - - No - The total number of DMAs that loaded this tensor into the State Buffer. * - ``load_to_sbuf_repeat_factor`` - number (float) - - No - The number of times this tensor was loaded into the State Buffer. A value of 1 means it was loaded once, 2 means it was loaded twice, and so on. .. _profile-schema-throttle: Throttle ~~~~~~~~ Hardware throttling notification events that track thermal or power throttling on NeuronCores. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``duration`` - integer (int64) - nanoseconds - Yes - Duration of the throttling event. * - ``k`` - integer (int32) - - Yes - Allowed utilization level value used to calculate throttling limit (utilization_limit = k/n, throttling = 1 - utilization_limit). * - ``n`` - integer (int32) - - Yes - Maximum possible utilization level value used to calculate throttling limit (utilization_limit = k/n, throttling = 1 - utilization_limit). * - ``timestamp`` - integer (int64) - nanoseconds - Yes - Timestamp when the throttling event occurred. * - ``type`` - string - - Yes - Type of throttling (e.g., "activity_0", "power_0", "temp_0"). A type prefixed with "unknown" means the throttle cause was not identified and its magnitude could not be measured; k and n are both reported as 1 in that case, making the utilization limit they imply incomplete. Example: ``power_0``. * - ``pcore_idx`` - integer (int64) - - No - Index of the physical NeuronCore within a Logical NeuronCore (LNC). A Logical NeuronCore is a group of physical NeuronCores. For LNC2 this field can have a value of either 0 or 1. .. _profile-schema-throttlesummary: ThrottleSummary ~~~~~~~~~~~~~~~ Summary statistics for hardware throttling on each NeuronCore and for each type of throttling. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``nc_idx`` - integer (int32) - - Yes - Index of the NeuronCore within a Neuron Device. * - ``throttle_type`` - string - - Yes - Type of throttling (e.g., "activity_0", "power_0", "temp_0"). The value "all" means the metric aggregates across all throttle types. A type prefixed with "unknown" means the throttle cause was not identified and its magnitude could not be measured. Example: ``power_0``. * - ``throttle_active_time_ns`` - integer (int64) - nanoseconds - No - Total time in nanoseconds when throttling was active. * - ``throttle_active_time_percent`` - number (float) - fraction (0-1) - No - Percent of total time when throttling was active. Range: 0 to 1. * - ``max_achievable_util_with_throttling_percent`` - number (float) - fraction (0-1) - No - Time-weighted average of the maximum achievable utilization with throttling, representing what fraction of compute capacity was permitted during throttled periods. A value of 1.0 normally means no throttling occurred, while 0.5 means the hardware was limited to 50% utilization. Calculated as the average of K/N ratios from HAM notifications weighted by the duration each limit was active. Throttling whose magnitude could not be measured has K and N both reported as 1, which also yields 1.0 here, making this average incomplete when throttle_active_time_percent is non-zero. Range: 0 to 1. .. _profile-schema-timelineannotation: TimelineAnnotation ~~~~~~~~~~~~~~~~~~ A time region flagged during profile processing as a potential performance or data-quality concern, rendered as a full-width highlight band across the timeline. Unlike the profile-level Warning table (which is non-positional), each row has a start/end and a type that the UI maps to a color and label. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``start_ts`` - integer (int64) - nanoseconds - Yes - Timestamp the flagged region started. * - ``end_ts`` - integer (int64) - nanoseconds - Yes - Timestamp the flagged region ended. * - ``duration_ns`` - integer (int64) - nanoseconds - Yes - Duration of the flagged region. * - ``type`` - string - - Yes - Machine-readable discriminator for the kind of region, used by the UI to pick a color and label. Extend as new region types are added. Example: ``small_dge_dma``. * - ``name`` - string - - Yes - Short human-readable label shown on the region. Example: ``Small DMA transfers``. * - ``description`` - string - - No - Longer explanation of the region and any remediation guidance, shown on hover. Example: ``Excessive small DMA transfers make instruction timing in this region unreliable.``. * - ``instruction_count`` - integer (int64) - - No - Number of contributing instructions merged into this region, when applicable. .. _profile-schema-warning: Warning ~~~~~~~ Warnings produced during profile processing. Used to communicate the possibility of missing or unexpected data. .. list-table:: :header-rows: 1 :widths: 22 16 10 8 44 * - Field - Type - Unit - Required - Description * - ``message`` - string - - Yes - The message displayed to users to describe the warning. Example: ``Current NEFF is missing compiler metrics. Please re-compile using latest version of the compiler.``. * - ``category`` - string - - Yes - The type of warning captured during profile processing. Example: ``Missing Data``. Enums ----- .. _profile-schema-capturetype: CaptureType ~~~~~~~~~~~ Where the profiled activity actually ran. 'hardware' is a capture from a Neuron device. 's0' and 'perf-sim' are simulated: 's0' is the instruction set simulator, 'perf-sim' is the compiler's post-scheduler performance simulation. Timings from a simulated capture are modeled, not measured, so they are not comparable to a hardware capture of the same model. Allowed values: * ``hardware`` * ``s0`` * ``perf-sim`` .. _profile-schema-datatype: DataType ~~~~~~~~ Canonical numeric data type of an instruction operand, normalized from the hardware ISA dtype enum (and, on the microscaling / compiler paths, from compiler-reported dtype strings). "unknown" is a fallback for any ISA code or dtype spelling that is not yet mapped to a canonical member; it makes newly introduced data types easy to detect so this enum can be extended. Operands with no applicable data type are omitted (null) rather than given a sentinel. Allowed values: * ``fp64`` * ``fp32`` * ``fp32r`` * ``tf32`` * ``fp16`` * ``bf16`` * ``fp8e3`` * ``fp8e4`` * ``fp8e5`` * ``mxfp8`` * ``mxfp4`` * ``int8`` * ``int16`` * ``int32`` * ``int64`` * ``uint8`` * ``uint16`` * ``uint32`` * ``uint64`` * ``unknown`` .. _profile-schema-dependencytype: DependencyType ~~~~~~~~~~~~~~ Kind of dependency edge in the Flow table. 'ORDERED', 'ANTI_DEPENDENCE', 'OUTPUT_DEPENDENCE' and 'FLOW_DEPENDENCE' come from the compiler's instruction debug info and describe a register or memory relationship between two instructions. The profiler adds three DMA kinds: 'trigger' is a PDMAT instruction that starts a DMA, 'completion' is an EVSEM instruction that waits for one, and 'DGE_TRIGGER' is a DGE instruction that starts a dynamic DMA. 'SAME_PSUM_BANK' links two instructions that used one PSUM bank on one core with no other access to that bank between them, where at least one of the two wrote the bank. The profiler derives 'SAME_PSUM_BANK' from decoded operand addresses and execution order rather than from a compiler or runtime record, so a pair of instructions can carry both a 'SAME_PSUM_BANK' edge and a compiler edge. 'INVALID' is the compiler's unset edge kind. A row whose dependency kind is unknown omits the field instead of naming a sentinel. Allowed values: * ``ORDERED`` * ``ANTI_DEPENDENCE`` * ``OUTPUT_DEPENDENCE`` * ``FLOW_DEPENDENCE`` * ``trigger`` * ``completion`` * ``DGE_TRIGGER`` * ``SAME_PSUM_BANK`` * ``INVALID`` .. _profile-schema-dmaqueuetype: DmaQueueType ~~~~~~~~~~~~ The type of a DMA queue. Identifies the origin and contents of DMA packets. ``software_dynamic`` packets are generated by the Gpsimd cores. ``hardware_dynamic`` packets are generated by dedicated DGE hardware. ``instruction``, ``input``, ``output``, and ``data`` packets are generated by the compiler. ``collectives`` packets move data for collective operations. Allowed values: * ``instruction`` * ``input`` * ``output`` * ``data`` * ``software_dynamic`` * ``hardware_dynamic`` * ``collectives`` * ``unknown`` .. _profile-schema-errortype: ErrorType ~~~~~~~~~ Category or type of error in the Error table. ``Errors`` are hardware error events. ``Warnings`` are hardware warning events. ``Custom Notification`` are custom notification events from the NeuronCore. Allowed values: * ``Custom Notification`` * ``Errors`` * ``Warnings`` .. _profile-schema-executionunit: ExecutionUnit ~~~~~~~~~~~~~ Classification of whether an instruction is executed by the sequencer (control plane) or dispatched to the compute engine (data plane). Allowed values: * ``compute`` * ``sequencer`` .. _profile-schema-memorybandwidthdirection: MemoryBandwidthDirection ~~~~~~~~~~~~~~~~~~~~~~~~ Direction of a memory bandwidth transfer between host and device. 'host_to_device' corresponds to dmem_buf_copyin events (system profile). 'device_to_host' corresponds to dmem_buf_copyout events (system profile). 'bidirectional' is the simultaneous sum of both directions and only appears as a computed series direction, never on an individual event. Allowed values: * ``host_to_device`` * ``device_to_host`` * ``bidirectional`` .. _profile-schema-performancemode: PerformanceMode ~~~~~~~~~~~~~~~ Performance mode for a Vector or Scalar Engine instruction. Indicates the throughput optimization applied by the hardware. Allowed values: * ``2x`` * ``4x`` .. _profile-schema-spilltype: SpillType ~~~~~~~~~ The spill/reload role of a DMA, when it is one. The compiler spills an intermediate tensor by saving it from State Buffer to an HBM scratch (temp) buffer and later reloading it back. 'save' is a State Buffer -> HBM temp transfer, 'reload' is HBM temp -> State Buffer. Derived from the DMA's source and destination memory-region types. Allowed values: * ``save`` * ``reload``