This document is relevant for: Inf1, Inf2, Trn1, Trn2, Trn3
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 Schema Reference
section below is generated directly from that definition, so it matches exactly
what neuron-explorer --show-profile-schema prints. This page reflects schema version 7.4.0 released in version 2.31.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:
neuron-explorer --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 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:
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). |
|
Time-series samples |
Each is a value at a point in time. |
|
Dependency-graph edges |
Each is a directed edge between rows in timeline tables (e.g. an instruction and the DMA it triggered). |
|
Hierarchical aggregation |
Each is a node in the compiler IR hierarchy (Framework → HLO → Penguin → BIR → Instruction), with rolled-up statistics. |
|
Aggregated summary |
Each is a summary metric computed across the whole profile (or by some key). |
|
Reference / lookup |
Each is a entity in the model or execution that other tables join to them by id to enrich rows without duplicating the detail. |
|
Profile metadata |
Each describes a piece of profile metadata. |
|
Diagnostics |
Each is a warning emitted during ingestion. Worth checking first, rows here can indicate a data-quality problem. |
|
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 and is viewable using the Database Viewer.
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 has 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
Enums
Tables#
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.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int64) |
nanoseconds |
Yes |
Timestamp when activity on this engine starts. |
|
integer (int64) |
nanoseconds |
Yes |
Timestamp when activity on this engine ends. |
|
string |
Yes |
Compute engine that this active time period corresponds to. Example: |
|
|
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. |
|
|
integer (int64) |
nanoseconds |
Yes |
The duration of the active time period. |
AssemblyInstruction#
Assembly instructions from the NEFF. These are lower level instructions than BIR instructions. Each BIR 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.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int64) |
Yes |
Foreign key reference to the parent Instruction.id. Links this assembly instruction to its parent BIR instruction. |
|
|
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: |
|
|
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: |
|
|
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. |
BirInstruction#
BIR level instruction information for displaying hierarchy graph. Parent level is penguin level, child level is instruction level.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int64) |
Yes |
Unique identifier for corresponding assembly level instruction. |
|
|
integer (int64) |
Yes |
Unique identifier for bir level instruction. |
|
|
integer (int64) |
Yes |
Unique identifier for penguin level instruction. |
|
|
integer (int64) |
Yes |
Unique identifier for hlo level instruction. |
|
|
integer (int64) |
Yes |
Unique identifier for framework level instruction. |
|
|
string |
Yes |
BIR instruction name. Example: |
|
|
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: |
|
|
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. |
|
|
integer (int64) |
nanoseconds |
Yes |
Minimum start time among all of the corresponding instruction timestamps. Unit is nanoseconds. |
|
integer (int64) |
nanoseconds |
Yes |
Maximum end time among all of the corresponding instruction timestamps. Unit is nanoseconds. |
|
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: |
|
|
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: |
|
|
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: |
|
|
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: |
|
integer (int64) |
bytes |
Yes |
Total number of bytes read from HBM by all grouped instructions in this BIR instruction. |
|
integer (int64) |
bytes |
Yes |
Total number of bytes written to HBM by all grouped instructions in this BIR instruction. |
|
integer (int64) |
bytes |
No |
Total number of bytes read from SBUF by all grouped instructions in this BIR instruction. |
|
integer (int64) |
bytes |
No |
Total number of bytes written to SBUF by all grouped instructions in this BIR instruction. |
|
integer (int64) |
bytes |
No |
Total number of bytes loaded by LDWEIGHTS instructions in this BIR instruction. |
CcInstruction#
Collective Communication (CC) instructions that are executed by the CC-Cores.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int64) |
nanoseconds |
Yes |
Timestamp when the collective instruction started execution. |
|
integer (int64) |
nanoseconds |
Yes |
Timestamp when the collective instruction finished execution. |
|
integer (int64) |
nanoseconds |
Yes |
Duration of the collective instruction. |
|
integer (int64) |
Yes |
The index of the CC-Core that executed the collective instruction. |
|
|
string |
Yes |
The name of the collective instruction as provided by NeuronCollectivesFirmware. Example: |
|
|
string |
No |
The engine that triggered the collective instruction. Example: |
|
|
integer (int64) |
nanoseconds |
No |
Start timestamp of the PSEUDO_TRIGGER_COLLECTIVE instruction which triggered this collective instruction. |
CcOp#
Collective Communication (CC) operations that are executed by the collective engines.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int64) |
nanoseconds |
Yes |
Timestamp when the collective operation started execution. |
|
integer (int64) |
nanoseconds |
Yes |
Timestamp when the collective operation finished execution. |
|
integer (int64) |
nanoseconds |
Yes |
Duration of the collective operation. |
|
string |
Yes |
The name of the collective operation being performed. Example: |
|
|
string |
Yes |
The algorithm used for performing the collective operation. Example: |
|
|
integer (int64) |
Yes |
Number of elements in the collectives operation. |
|
|
string |
Yes |
Data type used for the collective operation. Example: |
|
|
integer (int64) |
Yes |
ID of the collective operation. |
|
|
integer (int64) |
Yes |
Unique ID of the stream that the collective operation is on for a given profile. |
|
|
integer (int64) |
Yes |
ID of the stream that the collective operation is on relative to the execution. |
|
|
string |
Yes |
The engine that triggered the collective operation. Example: |
|
|
integer (int64) |
bytes |
No |
The input size in bytes of the collective operation. |
|
integer (int64) |
bytes |
No |
The output size in bytes of the collective operation. |
|
number (float) |
bytes_per_second |
No |
The algorithm bandwidth of the collective operation. Calculated as (max(input_size, output_size) / duration)). |
|
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 dependent on the operation and number of cores. |
|
integer (int64) |
nanoseconds |
No |
Start timestamp of the PSEUDO_TRIGGER_COLLECTIVE instruction which triggered this collective operation. |
|
integer (int64) |
nanoseconds |
No |
The time delay between when the collective operation was triggered, and when execution actually started. |
|
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. |
|
|
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: |
CcStream#
Information about Collective Communication (CC) streams.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int64) |
Yes |
The id of this collective communication stream. |
|
|
integer (int64) |
Yes |
ID of the stream that the collective operation is on relative to the execution. |
|
|
string |
Yes |
The CC-Core ids assigned to this stream. Example: |
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.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int64) |
nanoseconds |
Yes |
Timestamp the core barrier operation started. |
|
integer (int64) |
nanoseconds |
Yes |
Timestamp the core barrier operation completed. |
|
integer (int64) |
nanoseconds |
Yes |
Duration of the core barrier operation. |
|
string |
Yes |
The compute engine that this core barrier operation corresponds to. Example: |
|
|
integer (int64) |
Yes |
Unique identifier for this core barrier operation. |
|
|
integer (int64) |
Yes |
Semaphore used to implement this core barrier. |
|
|
integer (int64) |
Yes |
Compiler program counter for this core barrier instruction. |
|
|
string |
Yes |
Core barrier opcode at compile time. Example: |
|
|
string |
Yes |
Operands at compile time specifying semaphore conditions, completion events, and barrier ID. Example: |
|
|
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. |
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.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int64) |
nanoseconds |
Yes |
Timestamp for this CPU utilization measurement. |
|
number (float) |
percent |
Yes |
CPU utilization percentage at this timestamp. Each CPU utilization value represents the average CPU utilization over the previous 1 second time period. |
|
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: |
|
|
string |
Yes |
Identifier for the instance on which this CPU utilization was measured. Example: |
|
|
integer (int64) |
No |
Unique UltraServer ID if this measurement was captured on an UltraServer instance. |
|
|
integer (int64) |
No |
The instance index within the UltraServer on which this measurement was captured. |
DeviceProfileList#
This table contains device profile information including device profile names, instance IDs, process IDs, NeuronCore IDs, session IDs, and model IDs.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
string |
Yes |
Name of the device profile. Example: |
|
|
string |
Yes |
Unique identifier of the server instance that produced this profile. Example: |
|
|
integer (int64) |
Yes |
Process ID that produced this profile. |
|
|
integer (int64) |
Yes |
NeuronCore ID associated with this profile. |
|
|
integer (int64) |
No |
Unique session identifier for this profile. Present for session-based profiling. |
|
|
integer (int64) |
No |
Unique model identifier for this profile. Present for model-based profiling. |
DmaPacket#
A DMA packet transfers data within a Neuron Device. A DMA packet typically transfers a portion of a tensor. Source and destination locations can be HBM, SBUF, or PSUM. The DMA packet traverses from the source location, over the AXI bus, to a DMA engine, over the AXI bus, and to the destination location. DMA packets correspond to an aggregated DMA via the packet_aggregation_id field in the DmaPacket table which corresponds to an id the DmaPacketAggregated table. The corresponding row in the DmaPacketAggregated table contains more information about a packet such as the source, destination, variable name, and other values that are the same for all packets in the aggregated set of DMA packets.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int64) |
nanoseconds |
Yes |
Timestamp when instruction started execution. |
|
integer (int64) |
nanoseconds |
Yes |
Timestamp when instruction finished execution. |
|
integer (int64) |
nanoseconds |
Yes |
Duration of instruction execution. |
|
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. |
|
|
integer (int64) |
Yes |
Unique identifier for this DMA packet. |
|
|
integer (int32) |
Yes |
Index of the DMA packet within the set of DMA packets which fulfill a transfer of one or more tensors. |
|
|
integer (int64) |
Yes |
Unique identifier of the row in the DmaPacketAggregated table that this packet is aggregated in. Each aggregated DMA represents the transfer of one or more tensors. The transfer is fulfilled on a Neuron Device by a set of DMA packets. See the DmaPacketAggregated table for more details. The row in the DmaPacketAggregated table with this packet_aggregation_id has more information about this packet such as the source, destination, variable name, and other values that are the same for all packets in the aggregated set of DMA packets. |
|
|
integer (int64) |
No |
Index of the DMA engine that executed this DMA packet. |
|
|
integer (int64) |
bytes |
Yes |
Size in bytes of data read during this operation. |
|
integer (int64) |
bytes |
Yes |
Size in bytes of data written during this operation. |
|
integer (int64) |
bytes |
Yes |
Size in bytes of data transferred during this operation. |
|
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. |
|
string (DmaQueueType) |
Yes |
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: |
|
|
string |
No |
Variable or tensor name associated with the data being transferred. Example: |
|
|
string |
No |
Queue name where this DMA operation was queued. Example: |
|
|
integer (int64) |
No |
Queue index where this DMA operation was queued. |
|
|
integer (int64) |
Yes |
Cyclic Redundancy Check value for data integrity verification. Example: |
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 corresponds to the packet_aggregation_id field in the DmaPacket table, linking each DMA packet to its corresponding aggregated DMA.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int64) |
nanoseconds |
Yes |
Timestamp when instruction started execution. |
|
integer (int64) |
nanoseconds |
Yes |
Timestamp when instruction finished execution. |
|
integer (int64) |
nanoseconds |
Yes |
Duration of instruction execution. |
|
integer (int64) |
Yes |
Unique identifier for this aggregated set of DMA packets. |
|
|
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. |
|
|
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. |
|
|
integer (int64) |
Yes |
Queue index where this DMA operation was queued. |
|
|
string |
Yes |
Queue name where this DMA operation was queued. Example: |
|
|
string (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: |
|
|
string |
Yes |
DMA operation type. Example: |
|
|
string |
Yes |
Variable or tensor name associated with the data being transferred. Example: |
|
|
integer (int64) |
bytes |
Yes |
Size in bytes of data read during this operation. |
|
integer (int64) |
bytes |
Yes |
Size in bytes of data written during this operation. |
|
integer (int64) |
bytes |
Yes |
Size in bytes of data transferred during this operation. |
|
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. |
|
string |
Yes |
Source memory location(s) for the DMA transfer. Example: |
|
|
string |
Yes |
Destination memory location(s) for the DMA transfer. Example: |
|
|
string |
Yes |
Identifier of semaphore used for synchronization. Example: |
|
|
string |
No |
Modular flow function name and instance associated with this DMA operation. Example: |
|
|
string |
Yes |
State Buffer partitions read from during this DMA operation. Example: |
|
|
integer (int32) |
Yes |
Number of State Buffer partitions read from during this DMA operation. |
|
|
string |
Yes |
Multi-dimensional shape read from source memory. Example: |
|
|
string |
Yes |
Byte offset within the source memory region. Example: |
|
|
string |
Yes |
Stride of reads from source location. Example: |
|
|
string |
Yes |
State Buffer partitions written to during this DMA operation. Example: |
|
|
integer (int32) |
Yes |
Number of State Buffer partitions written to during this DMA operation. |
|
|
string |
Yes |
Multi-dimensional shape written to destination memory. Example: |
|
|
string |
Yes |
Byte offset within the destination memory region. Example: |
|
|
string |
Yes |
The destination stride. Example: |
|
|
boolean |
No |
Indicates if transpose mode is enabled. Example: |
|
|
string |
No |
Element size for transpose operations. Example: |
|
|
string |
No |
Shape descriptor for transpose operation. Example: |
|
|
string |
No |
Source shape for transpose operations. Example: |
|
|
string |
No |
Target shape for transpose operation. Example: |
|
|
string |
No |
Unroll factor for transpose operation. Example: |
|
|
integer (int64) |
No |
Start timestamp of the instruction that triggered this DMA. This DMA had a dependency on that instruction. |
|
|
integer (int64) |
No |
End timestamp of the instruction that triggered this DMA. This DMA had a dependency on that instruction. |
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.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
string |
Yes |
The name of the DMA queue. Example: |
|
|
string |
Yes |
The type of DMA queue. Example: |
|
|
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. |
|
|
integer (int64) |
Yes |
The index of this DMA queue within the DMA engine. |
|
|
integer (int64) |
Yes |
The index of the DMA engine that this DMA queue belongs to. |
|
|
integer (int64) |
Yes |
The AXI port used by this DMA queue. |
|
|
integer (int64) |
Yes |
The DRAM channel used by this DMA queue. |
DmaUsage#
DMA time-series usage metrics such as throughput over time.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int64) |
nanoseconds |
Yes |
Timestamp for this measurement. |
|
number (float) |
bytes_per_second |
Yes |
Rate of data transfer over DMA in bytes per second at this timestamp. |
|
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. |
|
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. |
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.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int64) |
nanoseconds |
Yes |
Bucket start timestamp aligned to bucket boundary. |
|
integer (int64) |
Yes |
Index of the physical NeuronCore. |
|
|
number (double) |
Yes |
Sum of current_throughput values in this bucket. Divide by count for average. |
|
|
integer (int64) |
Yes |
Sum of total_transfer_bytes values in this bucket. Divide by count for average. |
|
|
integer (int64) |
Yes |
Maximum total_transfer_bytes value in this bucket. |
|
|
integer (int64) |
Yes |
Number of raw rows aggregated into this bucket. |
|
|
integer (int64) |
nanoseconds |
Yes |
The bucket resolution in nanoseconds used during ingestion. |
Error#
Information about hardware generated error events such as Out Of Bounds (OOB) or Not A Number (NAN).
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
string |
Yes |
Short description of the error condition. Example: |
|
|
integer (int32) |
Yes |
Numerical value of error type. Distinguishes between different error conditions. Example: |
|
|
string |
No |
Error metadata as string. Example: |
|
|
string |
No |
Hexadecimal representation of error metadata. Example: |
|
|
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. |
|
|
integer (int64) |
nanoseconds |
Yes |
Timestamp when the error occurred. |
|
string (ErrorType) |
Yes |
Category or type of error. Allowed values: |
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.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int32) |
Yes |
Zero-based index of this execution within the profile. |
|
|
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: |
|
|
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. |
|
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. |
|
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. |
|
|
integer (int64) |
No |
Number of rows written to the profile database for this execution. |
Flow#
Storing dependency information for rows in different or same table (such as between instructions and DMAs).
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int64) |
Yes |
Unique identifier for the input flow. |
|
|
integer (int64) |
Yes |
Unique identifier for the output flow. |
|
|
string |
Yes |
Table name for the input flow. Example: |
|
|
string |
Yes |
Table name for the output flow. Example: |
|
|
string |
No |
Type of dependency from the predecessor instruction. Example: |
|
|
integer (int64) |
nanoseconds |
No |
The gap time between the end of the input event and the start of the output event. |
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.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int64) |
Yes |
Unique identifier for corresponding assembly level instruction. |
|
|
integer (int64) |
Yes |
Unique identifier for bir level instruction. |
|
|
integer (int64) |
Yes |
Unique identifier for penguin level instruction. |
|
|
integer (int64) |
Yes |
Unique identifier for hlo level instruction. |
|
|
integer (int64) |
Yes |
Unique identifier for framework level instruction. |
|
|
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: |
|
|
string |
No |
Framework operation name. Example: |
|
|
integer (int64) |
nanoseconds |
Yes |
Minimum start time among all of the corresponding instruction timestamps. Unit is nanoseconds. |
|
integer (int64) |
nanoseconds |
Yes |
Maximum end time among all of the corresponding instruction timestamps. Unit is nanoseconds. |
|
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: |
|
|
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: |
|
|
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: |
|
|
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: |
|
integer (int64) |
bytes |
Yes |
Total number of bytes read from HBM by all grouped instructions in this Framework instruction. |
|
integer (int64) |
bytes |
Yes |
Total number of bytes written to HBM by all grouped instructions in this Framework instruction. |
|
integer (int64) |
bytes |
No |
Total number of bytes read from SBUF by all grouped instructions in this Framework instruction. |
|
integer (int64) |
bytes |
No |
Total number of bytes written to SBUF by all grouped instructions in this Framework instruction. |
|
integer (int64) |
bytes |
No |
Total number of bytes loaded by LDWEIGHTS instructions in this Framework instruction. |
|
integer (int64) |
No |
Total number of instructions executed by PE (Tensor) engine in this Framework instruction. |
|
|
integer (int64) |
No |
Total number of instructions executed by ACT (Scalar) engine in this Framework instruction. |
|
|
integer (int64) |
No |
Total number of instructions executed by SP (Sync) engine in this Framework instruction. |
|
|
integer (int64) |
No |
Total number of instructions executed by POOL (GpSimd) engine in this Framework instruction. |
|
|
integer (int64) |
No |
Total number of instructions executed by DVE (Vector) engine in this Framework instruction. |
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.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
string |
Yes |
Name of the current node in the hierarchy. Example: |
|
|
string |
Yes |
Comma separated array of child node names. Example: |
|
|
string |
Yes |
Name of the parent node in the hierarchy. Example: |
|
|
string |
Yes |
Comma separated array of framework operator IDs associated with this node. Example: |
|
|
integer (int64) |
nanoseconds |
Yes |
Earliest start time among all operations in this node. Unit is nanoseconds. |
|
integer (int64) |
nanoseconds |
Yes |
Latest end time among all operations in this node. Unit is nanoseconds. |
|
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: |
|
|
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: |
|
|
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: |
|
|
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: |
|
number (float) |
percent |
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)). |
|
number (float) |
percent |
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)). |
|
string |
No |
Comma separated array of per-interval Hardware FLOPs Utilization (HFU) percentages corresponding to the time intervals in time_intervals field. Example: |
|
|
string |
No |
Comma separated array of per-interval Model FLOPs Utilization (MFU) percentages corresponding to the time intervals in time_intervals field. Example: |
|
|
string |
No |
Comma separated array of per-interval Memory Bandwidth Utilization (MBU) percentages corresponding to the time intervals in time_intervals field. Example: |
|
|
integer (int64) |
bytes |
No |
Total number of bytes read from HBM by all grouped instructions in this framework node. |
|
integer (int64) |
bytes |
No |
Total number of bytes written to HBM by all grouped instructions in this framework node. |
|
integer (int64) |
bytes |
No |
Total number of bytes read from SBUF by all grouped instructions in this framework node. |
|
integer (int64) |
bytes |
No |
Total number of bytes written to SBUF by all grouped instructions in this framework node. |
|
integer (int64) |
No |
Total number of instructions executed by Tensor engine in this framework node. |
|
|
integer (int64) |
No |
Total number of instructions executed by Scalar engine in this framework node. |
|
|
integer (int64) |
No |
Total number of instructions executed by Sync engine in this framework node. |
|
|
integer (int64) |
No |
Total number of instructions executed by GpSimd engine in this framework node. |
|
|
integer (int64) |
No |
Total number of instructions executed by Vector engine in this framework node. |
|
|
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: |
|
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: |
|
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: |
|
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: |
|
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: |
|
integer (int64) |
bytes |
No |
Total number of bytes loaded by LDWEIGHTS instructions in this framework node. |
|
integer (int64) |
No |
Total number of channels processed by Activation (Scalar) engine instructions in this framework node. |
|
|
integer (int64) |
No |
Total number of channels processed by DVE (Vector) engine instructions in this framework node. |
HbmUsage#
HBM time-series usage metrics such as throughput over time.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int64) |
nanoseconds |
Yes |
Timestamp for this measurement. |
|
number (float) |
bytes_per_second |
Yes |
Rate of data transfer to/from HBM in bytes per second at this point of time. |
|
integer (int64) |
bytes |
Yes |
Total bytes transferred to/from HBM. This total is cumulative from the start of profiling to this timestamp. |
|
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. |
HbmUsageSummaryByType#
HBM memory usage information summarized by type.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int64) |
Yes |
The Neuron Device index. |
|
|
integer (int64) |
Yes |
The NeuronCore index. |
|
|
string |
Yes |
The type of memory usage. Example: |
|
|
integer (int64) |
bytes |
Yes |
The amount of HBM used by this type of usage. |
HloInstruction#
HLO level instruction information for displaying hierarchy graph. Parent level is framework level, child level is penguin level.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int64) |
Yes |
Unique identifier for corresponding assembly level instruction. |
|
|
integer (int64) |
Yes |
Unique identifier for bir level instruction. |
|
|
integer (int64) |
Yes |
Unique identifier for penguin level instruction. |
|
|
integer (int64) |
Yes |
Unique identifier for hlo level instruction. |
|
|
integer (int64) |
Yes |
Unique identifier for framework level instruction. |
|
|
string |
Yes |
High-Level Optimizer (HLO) operation name. Example: |
|
|
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: |
|
|
integer (int64) |
nanoseconds |
Yes |
Minimum start time among all of the corresponding instruction timestamps. Unit is nanoseconds. |
|
integer (int64) |
nanoseconds |
Yes |
Maximum end time among all of the corresponding instruction timestamps. Unit is nanoseconds. |
|
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: |
|
|
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: |
|
|
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: |
|
|
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: |
|
integer (int64) |
bytes |
Yes |
Total number of bytes read from HBM by all grouped instructions in this HLO instruction. |
|
integer (int64) |
bytes |
Yes |
Total number of bytes written to HBM by all grouped instructions in this HLO instruction. |
|
integer (int64) |
bytes |
No |
Total number of bytes read from SBUF by all grouped instructions in this HLO instruction. |
|
integer (int64) |
bytes |
No |
Total number of bytes written to SBUF by all grouped instructions in this HLO instruction. |
|
integer (int64) |
bytes |
No |
Total number of bytes loaded by LDWEIGHTS instructions in this HLO instruction. |
HostMemUsage#
Host memory utilization time-series metrics. Each row represents a memory usage sample for an instance.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int64) |
nanoseconds |
Yes |
Timestamp for this memory utilization measurement. |
|
integer (int64) |
bytes |
Yes |
Amount of host memory currently in use. |
|
number (float) |
percent |
Yes |
Host memory utilization percentage at this timestamp. Calculated as (used_bytes / memory_capacity) * 100. |
|
string |
Yes |
Identifier for the instance on which this memory utilization was measured. Example: |
|
|
integer (int64) |
No |
Unique UltraServer ID if this measurement was captured on an UltraServer instance. |
|
|
integer (int64) |
No |
The instance index within the UltraServer on which this measurement was captured. |
Instruction#
Instructions executed on NeuronCore compute engines.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int64) |
nanoseconds |
Yes |
Timestamp when instruction started execution. |
|
integer (int64) |
nanoseconds |
Yes |
Timestamp when instruction finished execution. |
|
integer (int64) |
nanoseconds |
Yes |
Duration of instruction execution. |
|
string |
Yes |
Compute engine that executed this instruction. Example: |
|
|
string |
No |
The name of the timeline track that this instruction appears on in the Neuron Explorer web UI. Example: |
|
|
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. |
|
|
string |
Yes |
Opcode for the instruction. Example: |
|
|
string |
Yes |
Input operands for the instruction. Example: |
|
|
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 “comiler_opcode: PSEUDO_DMA_TRIGGER” may change to “opcode: WRITE” at runtime because that is when the write address is known. Example: |
|
|
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: |
|
|
integer (int64) |
Yes |
Unique identifier for this instruction execution. |
|
|
integer (int64) |
Yes |
The runtime program counter is this instruction index on the given engine. |
|
|
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. |
|
|
string |
No |
Penguin compiler ID containing both raw ID and context information. Example: |
|
|
integer (int64) |
No |
Raw internal penguin compiler identifier. |
|
|
string |
Yes |
Binary Intermediate Representation instruction identifier. Example: |
|
|
integer (int64) |
No |
Unprocessed Binary Intermediate Representation 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. |
|
|
string |
No |
Name of the BIR instruction from the name in the BIR debug info. Example: |
|
|
string |
No |
Unique kernel instruction identifier formatted as ‘bir_instruction_name:parent_tensorizer_id’, where parent_tensorizer_id is the parent penguin instruction ID for this instruction. Corresponds to the ‘kernel_instruction_name’ field in the KernelInstructions table. Example: |
|
|
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. |
|
|
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. |
|
|
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. |
|
string |
No |
The debug info associated with the HLO instruction that produced this instruction. Example: |
|
|
string |
No |
High-Level Optimizer (HLO) operation name. Example: |
|
|
integer (int64) |
No |
Unique identifier for hlo level instruction. |
|
|
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. |
|
|
string |
No |
Here “layer” means the name from the framework level debug info. Example: |
|
|
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: |
|
|
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: |
|
|
string |
No |
Source code location recorded in BIR debug info. This field can reference various source file types including NKI files, non-NKI Python files, and C++ files. Example: |
|
|
integer (int64) |
No |
Start timestamp of the DMA that this instruction triggered. |
|
|
integer (int64) |
No |
End timestamp of the DMA that this instruction triggered. |
|
|
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. |
|
string |
No |
List of instructions in this chain. Some instructions must be executed together in a sequential order such as custom operations or pseudo 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: |
|
|
integer (int64) |
No |
Number of bytes read from HBM by this instruction. |
|
|
integer (int64) |
No |
Number of bytes written to HBM by this instruction. |
|
|
integer (int64) |
No |
Number of bytes read from SBUF by this instruction. |
|
|
integer (int64) |
No |
Number of bytes written to SBUF by this instruction. |
|
|
integer (int64) |
No |
Number of bytes saved to HBM from SBUF by this instruction due to spill (lack of space in SBUF). |
|
|
integer (int64) |
No |
Number of bytes loaded from HBM to SBUF by this instruction due to spill (lack of space in SBUF). |
|
|
integer (int64) |
No |
Number of bytes read from PSUM by this instruction. |
|
|
integer (int64) |
No |
Number of bytes read from PSUM by this instruction. |
|
|
integer (int64) |
No |
Number of bytes read from PSUM into SBUF by this instruction. |
|
|
integer (int64) |
No |
Number of bytes loaded from an INPUT DMA queue by this instruction. |
|
|
integer (int64) |
No |
Number of bytes loaded from an OUTPUT DMA queue by this instruction. |
|
|
integer (int64) |
No |
Number of bytes loaded from a DATA / WEIGHT DMA queue by this instruction. |
|
|
string (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: |
|
|
string |
No |
Instruction type name from the debug info. Example: |
|
|
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. |
|
|
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. |
|
|
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. FP8 uses 1.0x on Trn1/Inf2 and 0.5x on Trn2/Trn3. UINT8 uses 0.5x. MXFP8 and MXFP4 use 0.25x on Trn3. |
|
|
integer (int64) |
No |
Number of data elements processed by this instruction, only set when tensor_instruction_type is TRANSPOSE. 0 for non-transpose instructions. |
|
|
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. |
|
|
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. BF16, FP16, and TF32 use a 1.0x multiplier. FP32 uses 4.0x. FP8 uses 1.0x on Trn1/Inf2 and 0.5x on Trn2/Trn3. UINT8 uses 0.5x. MXFP8 and MXFP4 use 0.25x on Trn3. |
|
|
string |
No |
Type of activation function used. Example: |
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).
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
string |
Yes |
Unique kernel instruction identifier formatted as ‘instruction_name:parent_tensorizer_id’, where parent_tensorizer_id is the parent penguin instruction ID for this instruction. parent_tensorizer_id refers to the parent. Example: |
|
|
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: |
|
|
integer (int64) |
No |
Total depth of the call stack (number of frames in the stack). Example: |
|
|
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: |
|
|
integer (int64) |
No |
Total depth of iteration variable nesting (number of nested loops). Example: |
|
|
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. |
|
|
string |
No |
Source file path where the code that generates this instruction is located. Example: |
|
|
integer (int64) |
No |
Line number in the source file that generates this instruction. Example: |
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).
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int64) |
Yes |
Unique identifier for this iteration variable entry. Example: |
|
|
string |
No |
Name of the iteration variable. Example: |
|
|
integer (int64) |
No |
Value of the iteration variable for this specific iteration. Example: |
|
|
integer (int64) |
No |
Reference to the outer loop’s iteration variable. Null for the outermost loop. Example: |
|
|
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. |
|
|
string |
No |
Source file path where this iteration variable is defined. Example: |
|
|
integer (int64) |
No |
Line number in the source file where this iteration variable is defined. Example: |
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).
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int64) |
Yes |
Unique identifier for this stack frame entry. Example: |
|
|
string |
No |
Fully qualified function name for this stack frame. Example: |
|
|
integer (int64) |
No |
Reference to the parent (caller) stack frame in the call stack. Null for the outermost function. Example: |
|
|
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. |
|
|
string |
No |
Source file path where this function is defined. Example: |
|
|
integer (int64) |
No |
Line number in the source file where this function is defined. Example: |
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.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int64) |
nanoseconds |
Yes |
Timestamp for this bandwidth measurement. Matches system profile trace event timestamps. |
|
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. |
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.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
array of MemoryBandwidthSeries |
Yes |
List of bandwidth time-series. Each series represents a unique combination of grouping key and transfer direction. Example: |
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.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
string |
Yes |
Identifier for the instance whose events are aggregated in this series. Example: |
|
|
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. |
|
|
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). |
|
|
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: |
||
|
array of MemoryBandwidthPoint |
Yes |
Time-ordered bandwidth samples. Each point represents the aggregate bandwidth at a unique timestamp where an event starts or ends. Example: |
Metadata#
Metadata associated with this profile, including details about the system configuration. Also see the NeffHeader table for metadata included in the NEFF header.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
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: |
|
|
integer (int32) |
Yes |
Version of the NTFF format used. This is the version of the raw profile data before processing. |
|
|
integer (int32) |
Yes |
The number of physical NeuronCores in the Logical NeuronCore. |
|
|
boolean |
Yes |
Whether execution was on Logical NeuronCores consisting of multiple physical cores. Example: |
|
|
integer (int64) |
nanoseconds |
Yes |
First timestamp in the trace. |
|
integer (int64) |
nanoseconds |
Yes |
Last timestamp in the trace. |
|
string |
Yes |
First timestamp in Unix time format. Example: |
|
|
string |
Yes |
Last timestamp in Unix time format. Example: |
|
|
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. |
|
integer (int32) |
Yes |
Number of DMA engines per physical NeuronCore. |
|
|
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. |
|
integer (int64) |
Yes |
The number of rows in the systolic array of the Tensor Engine. |
|
|
integer (int64) |
Yes |
The number of columns in the systolic array of the Tensor Engine. |
|
|
number (float) |
gigahertz |
Yes |
The clock frequency of the Tensor Engine. |
|
number (float) |
gigahertz |
Yes |
The clock frequency of the Vector Engine. |
|
integer (int64) |
Yes |
The number of channels in the Vector Engine. |
|
|
number (float) |
gigahertz |
Yes |
The clock frequency of the Scalar Engine. |
|
integer (int64) |
Yes |
The number of channels in the Scalar Engine. |
|
|
number (float) |
gigahertz |
Yes |
The clock frequency of the GpSimd Engine. |
|
integer (int64) |
Yes |
The number of channels in the GpSimd Engine. |
|
|
number (float) |
gigahertz |
Yes |
The clock frequency of the NeuronCore. |
|
integer (int64) |
Yes |
The number of ticks per nanosecond that the profile captures. Used to convert internal profile timestamps to nanoseconds. |
|
|
string |
No |
The instance type on which this profile was captured. Example: |
|
|
string |
No |
The Neuron Runtime library version with which this profile was captured. Example: |
|
|
string |
No |
The Neuron Collectives library version with which this profile was captured. Example: |
|
|
string |
No |
The Neuron driver version with which this profile was captured. Example: |
|
|
string |
No |
The Neuron Explorer version which post-processed this captured profile. Example: |
|
|
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: |
|
|
integer (int64) |
No |
The index of this NeuronCore. |
|
|
integer (int64) |
No |
The index of this NeuronDevice. |
|
|
integer (int64) |
No |
Unique UltraServer ID if this profile was captured on an UltraServer instance. |
|
|
string |
No |
The type of instances connectivity within an UltraServer. Example: |
|
|
integer (int32) |
No |
The total numbers of instance in an UltraServer. |
|
|
integer (int32) |
No |
The instance index within the UltraServer on which this profile was captured. |
|
|
string |
No |
Full S3 path where profile files (NEFF/NTFF/source/debug) are stored and uploaded. Example: |
|
|
string |
No |
User chosen name for this profile. Example: |
|
|
string |
No |
The date and time when the `neuron-explorer` tool was built. Example: |
|
|
boolean |
No |
Whether this profile is a post-scheduler simulation produced by the compiler. Example: |
NeffHeader#
Information contained in the NEFF header.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
string |
No |
Name of the NEFF. Example: |
|
|
string |
No |
Unique identifier of the NEFF. Example: |
|
|
string |
No |
Hash of the NEFF. Example: |
|
|
string |
No |
Comma-separated list of enabled features. Example: |
|
|
integer (int64) |
bytes |
No |
Size of the NEFF header in bytes. |
|
integer (int64) |
bytes |
No |
Data size in bytes. |
|
integer (int64) |
No |
NEFF major version number. |
|
|
integer (int64) |
No |
NEFF minor version number. |
|
|
string |
No |
Package version. Example: |
|
|
string |
No |
Build version. Example: |
|
|
integer (int64) |
No |
Number of collective communication (CC) participants involved. |
|
|
integer (int64) |
No |
Number of NeuronCores used by the NEFF. |
|
|
integer (int64) |
No |
Number of Logical NeuronCores used by the NEFF. |
|
|
integer (int64) |
No |
Number of physical NeuronCores per Logical NeuronCore (LNC). |
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.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
string |
Yes |
Compute engine that executed the instructions. Example: |
|
|
string |
Yes |
Opcode of the instructions. Example: |
|
|
integer (int64) |
Yes |
Total number of instructions with this engine and opcode combination. |
|
|
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). |
|
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. |
|
number (float) |
Yes |
Percentage of total execution time when at least one instruction with this engine and opcode combination is executing. |
PendingDma#
Outstanding DMA operations that are queued.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int64) |
nanoseconds |
Yes |
Time of pending DMA measurement. |
|
integer (int32) |
Yes |
Number of pending DMA operations queued. |
|
|
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. |
PenguinInstruction#
Penguin level instruction information for displaying hierarchy graph. Parent level is hlo level, child level is bir level.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int64) |
Yes |
Unique identifier for corresponding assembly level instruction. |
|
|
integer (int64) |
Yes |
Unique identifier for bir level instruction. |
|
|
integer (int64) |
Yes |
Unique identifier for penguin level instruction. |
|
|
integer (int64) |
Yes |
Unique identifier for hlo level instruction. |
|
|
integer (int64) |
Yes |
Unique identifier for framework level instruction. |
|
|
string |
Yes |
Penguin operation name. Example: |
|
|
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: |
|
|
integer (int64) |
nanoseconds |
Yes |
Minimum start time among all of the corresponding instruction timestamps. Unit is nanoseconds. |
|
integer (int64) |
nanoseconds |
Yes |
Maximum end time among all of the corresponding instruction timestamps. Unit is nanoseconds. |
|
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: |
|
|
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: |
|
|
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: |
|
|
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: |
|
integer (int64) |
bytes |
Yes |
Total number of bytes read from HBM by all grouped instructions in this Penguin instruction. |
|
integer (int64) |
bytes |
Yes |
Total number of bytes written to HBM by all grouped instructions in this Penguin instruction. |
|
integer (int64) |
bytes |
No |
Total number of bytes read from SBUF by all grouped instructions in this Penguin instruction. |
|
integer (int64) |
bytes |
No |
Total number of bytes written to SBUF by all grouped instructions in this Penguin instruction. |
|
integer (int64) |
bytes |
No |
Total number of bytes loaded by LDWEIGHTS instructions in this Penguin instruction. |
PsumUsage#
PSUM memory bank utilization metrics during execution.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int64) |
nanoseconds |
Yes |
Timestamp for this usage measurement. |
|
number (float) |
percent |
No |
Average fullness across all PSUM memory banks. |
|
number (float) |
percent |
No |
Fullness percentage of PSUM memory bank 0. |
|
number (float) |
percent |
No |
Fullness percentage of PSUM memory bank 1. |
|
number (float) |
percent |
No |
Fullness percentage of PSUM memory bank 2. |
|
number (float) |
percent |
No |
Fullness percentage of PSUM memory bank 3. |
|
number (float) |
percent |
No |
Fullness percentage of PSUM memory bank 4. |
|
number (float) |
percent |
No |
Fullness percentage of PSUM memory bank 5. |
|
number (float) |
percent |
No |
Fullness percentage of PSUM memory bank 6. |
|
number (float) |
percent |
No |
Fullness percentage of PSUM memory bank 7. |
|
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. |
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.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int64) |
nanoseconds |
Yes |
Timestamp when the memory address range becomes occupied (allocated). |
|
integer (int64) |
nanoseconds |
Yes |
Timestamp when the memory address range becomes free (deallocated). |
|
integer (int64) |
bytes |
Yes |
Starting memory address of the occupied region (inclusive - first byte of allocation). |
|
integer (int64) |
bytes |
Yes |
Ending memory address of the occupied region (inclusive - last byte of allocation). |
|
integer (int64) |
Yes |
State Buffer partition number that this address range belongs to (0, 32, 64, or 96). |
|
|
string |
No |
Source of memory operation. Distinguishes whether the memory allocation originated from a compute engine instruction or a DMA operation. Example: |
|
|
string |
No |
Instruction opcode name (only for instruction source type). Examples include MATMUL, TENSOR_TENSOR, etc. Example: |
|
|
string |
No |
Instruction operands string (only for instruction source type). Contains source/destination addresses and other instruction parameters. Example: |
|
|
string |
No |
DMA queue name (only for DMA source type). Identifies which DMA queue performed the operation. Example: |
|
|
integer (int64) |
No |
DMA block ID (only for DMA source type). Identifies which block within the DMA queue this operation corresponds to. |
|
|
string |
No |
Name of the tensor or variable associated with this memory allocation (if available). Example: |
|
|
string |
No |
High-Level Optimizer (HLO) operation name associated with this memory allocation (if available). Example: |
|
|
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. |
SbufUsage#
State Buffer (SB) memory partition utilization metrics during execution.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int64) |
nanoseconds |
Yes |
Timestamp for this usage measurement. |
|
number (float) |
percent |
No |
Average utilization across all State Buffer partitions. |
|
number (float) |
percent |
No |
Fullness percentage of State Buffer partition 0. |
|
number (float) |
percent |
No |
Fullness percentage of State Buffer partition 32. |
|
number (float) |
percent |
No |
Fullness percentage of State Buffer partition 64. |
|
number (float) |
percent |
No |
Fullness percentage of State Buffer partition 96. |
|
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. |
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.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
string |
Yes |
The table this field belongs to. Example: |
|
|
string |
Yes |
The description of the table this field belongs to. Example: |
|
|
string |
Yes |
The name of this field. Example: |
|
|
string |
Yes |
The description of this field. Example: |
|
|
string |
No |
One or more examples of what value this field may hold. Example: |
|
|
string |
Yes |
The data type of this field. Example: |
|
|
boolean |
Yes |
Indicates whether this field is required according to the OpenAPI schema definition. Example: |
|
|
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: |
|
|
string |
No |
The unit of this field. Example: |
|
|
number (float) |
No |
The minimum value that can be held by this field. Example: |
|
|
number (float) |
No |
The maximum value that can be held by this field. Example: |
|
|
string |
No |
Comma-separated list of allowed enum values for this field, if the field is an enum type. Example: |
SemaphoreUpdate#
Semaphore state changes used for synchronization between operations.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int64) |
nanoseconds |
Yes |
Timestamp when the semaphore update occurred. |
|
string |
Yes |
Unique identifier of the semaphore being updated. Example: |
|
|
integer (int32) |
Yes |
New value assigned to the semaphore. |
|
|
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: |
|
|
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: |
|
|
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: |
|
|
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: |
|
|
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: |
|
|
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. |
StackFrame#
Individual stack frames representing the call stack hierarchy for debugging and tracing execution back to source code.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
string |
Yes |
Unique identifier for this stack frame (one-based index converted to string). Example: |
|
|
integer (int32) |
Yes |
Reference to the file location ID in StackFrameFileLocation table. |
|
|
string |
Yes |
Reference to the parent stack frame ID, or 0 if this is the root frame. Example: |
StackFrameFileLocation#
File location information combining file names, function names, and line numbers for debugging and tracing execution back to source code.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
string |
Yes |
Unique identifier for this file location (one-based index converted to string). Example: |
|
|
integer (int32) |
Yes |
Reference to the file name ID in StackFrameFileName table. |
|
|
integer (int32) |
Yes |
Reference to the function name ID in StackFrameFunctionName table. |
|
|
integer (int32) |
Yes |
Line number within the source file. |
StackFrameFileName#
Source code file names referenced in the stack frames for debugging and tracing execution back to source code.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
string |
Yes |
Unique identifier for this file name (one-based index converted to string). Example: |
|
|
string |
Yes |
Source code file name or path. Example: |
StackFrameFunctionName#
Function names referenced in the stack frames for debugging and tracing execution back to source code.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
string |
Yes |
Unique identifier for this function name (one-based index converted to string). Example: |
|
|
string |
Yes |
Function or method name. Example: |
Summary#
Profile performance metrics which summarize execution. For profile metadata and system information see the Metadata table.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
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. |
|
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. |
|
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. |
|
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. |
|
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. |
|
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. FP8 uses 1.0x on Trn1/Inf2 and 0.5x on Trn2/Trn3. UINT8 uses 0.5x. MXFP8 and MXFP4 use 0.25x on Trn3. |
|
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. BF16, FP16, and TF32 use a 1.0x multiplier. FP32 uses 4.0x. FP8 uses 1.0x on Trn1/Inf2 and 0.5x on Trn2/Trn3. UINT8 uses 0.5x. MXFP8 and MXFP4 use 0.25x on Trn3. |
|
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). |
|
|
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. |
|
|
integer (int64) |
No |
Total number of trace events in the NTFF including instructions, events/semaphores, and DMAs. |
|
|
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. |
|
|
integer (int64) |
No |
Total number of data points written to the database. |
|
|
integer (int64) |
cycles |
No |
dma_active_time multiplied by TPB clock speed. |
|
number (float) |
seconds |
No |
Duration of time when at least one DMA packet is transferring. |
|
number (float) |
percent |
No |
Percentage of time when at least one DMA packet is transferring. |
|
number (float) |
seconds |
No |
Duration of time when at least one statically generated DMA packet is transferring. |
|
number (float) |
percent |
No |
Percentage of time when at least one dynamically generated DMA packet is transferring. |
|
number (float) |
percent |
No |
Percentage of time when at least one statically generated DMA packet is transferring. |
|
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. |
|
number (float) |
seconds |
No |
Total duration of all DMA engine packets. Overlapping DMA packets across all DMA engines are added together. |
|
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. |
|
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. |
|
|
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. |
|
integer (int64) |
No |
Total number of DMA queues specified in the NEFF including all input, output, and weight DMA queues for different DMA engines. |
|
|
integer (int64) |
bytes |
No |
Size of input tensors and weight tensors. |
|
integer (int64) |
bytes |
No |
Size of input tensors, output tensors, and weight tensors. Used to support perf bounds, specifically MemoryBoundIdealNoReloads |
|
integer (int64) |
bytes |
No |
Total size of all weights defined in the NEFF. |
|
integer (int64) |
No |
Number of DMA packets that are statically generated. |
|
|
integer (int64) |
bytes |
No |
Total transfer size of DMA packets that are statically generated. |
|
number (float) |
percent |
No |
Percentage of DMA packets that are dynamically generated during execution (generally higher means better performance). |
|
number (float) |
percent |
No |
Percentage of DMA packets that are statically generated. |
|
number (float) |
percent |
No |
Percentage of total transfer size for DMA packets that are dynamically generated during execution. |
|
number (float) |
percent |
No |
Percentage of total transfer size for DMA packets that are statically generated. |
|
number (float) |
percent |
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. |
|
number (float) |
percent |
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. |
|
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. |
|
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. |
|
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. |
|
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. |
|
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. |
|
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. |
|
integer (int64) |
bytes |
No |
Total bytes of data that are written to PSUM by compute engine instructions. |
|
integer (int64) |
bytes |
No |
Total bytes of data that are read from PSUM by compute engine instructions. |
|
integer (int64) |
No |
Number of compute engine instructions that moved data from PSUM to SBUF. |
|
|
integer (int64) |
bytes |
No |
Total bytes of data moved from PSUM to SBUF by compute engine instructions. |
|
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. |
|
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. |
|
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. |
|
number (float) |
seconds |
No |
Duration of time when at least one software generated DMA packet is transferring. |
|
number (float) |
percent |
No |
Percentage of time when at least one software generated DMA packet is transferring. |
|
integer (int64) |
No |
Number of DMA packets that are dynamically generated during execution via software backend. |
|
|
integer (int64) |
bytes |
No |
Total transfer size of DMA packets that are dynamically generated during execution via software backend. |
|
number (float) |
percent |
No |
Percentage of DMA packets that are dynamically generated during execution via software backend. |
|
number (float) |
percent |
No |
Percentage of total transfer size for DMA packets that are dynamically generated via software backend. |
|
number (float) |
seconds |
No |
Duration of time when at least one hardware generated DMA packet is transferring. |
|
number (float) |
percent |
No |
Percentage of time when at least one hardware generated DMA packet is transferring. |
|
integer (int64) |
No |
Number of DMA packets that are dynamically generated during execution via dedicated hardware. |
|
|
integer (int64) |
bytes |
No |
Total transfer size of DMA packets that are dynamically generated during execution via dedicated hardware. |
|
number (float) |
percent |
No |
Percentage of DMA packets that are dynamically generated during execution via dedicated hardware. |
|
number (float) |
percent |
No |
Percentage of total transfer size for DMA packets that are dynamically generated via dedicated hardware. |
|
number (float) |
seconds |
No |
Total duration of all Tensor engine instructions. Overlapping instructions on the engine are added together. Units are in seconds. |
|
number (float) |
seconds |
No |
Duration of time when Tensor engine is processing at least one instruction (excluding semaphore waits). Units are in seconds. |
|
number (float) |
percent |
No |
Percentage of time when Tensor engine is processing at least one instruction (excluding semaphore waits). |
|
integer (int64) |
No |
Number of Tensor Engine instructions. |
|
|
integer (int64) |
No |
Total number of MATMUL instructions. |
|
|
number (float) |
percent |
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. |
|
number (float) |
percent |
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. |
|
number (float) |
percent |
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. |
|
number (float) |
percent |
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. |
|
number (float) |
percent |
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. |
|
number (float) |
percent |
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. |
|
number (float) |
percent |
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. |
|
number (float) |
seconds |
No |
Total duration of all Vector engine instructions. Overlapping instructions on the engine are added together. Units are in seconds. |
|
number (float) |
seconds |
No |
Duration of time when Vector engine is processing at least one instruction (excluding semaphore waits). Units are in seconds. |
|
number (float) |
percent |
No |
Percentage of time when Vector engine is processing at least one instruction (excluding semaphore waits). |
|
integer (int64) |
No |
Number of Vector Engine instructions. |
|
|
number (float) |
seconds |
No |
Total duration of all Scalar engine instructions. Overlapping instructions on the engine are added together. Units are in seconds. |
|
number (float) |
seconds |
No |
Duration of time when Scalar engine is processing at least one instruction (excluding semaphore waits). Units are in seconds. |
|
number (float) |
percent |
No |
Percentage of time when Scalar engine is processing at least one instruction (excluding semaphore waits). |
|
integer (int64) |
No |
Number of Scalar Engine instructions. |
|
|
number (float) |
seconds |
No |
Total duration of all ACTIVATE/ACTIVATE_QUANTIZE instructions. Overlapping instructions are added up. Units are in seconds. |
|
integer (int64) |
No |
Number of ACTIVATE/ACTIVATE_QUANTIZE instructions. |
|
|
number (float) |
seconds |
No |
Total duration of all Sync engine instructions. Overlapping instructions on the engine are added together. Units are in seconds. |
|
number (float) |
seconds |
No |
Duration of time when Sync engine is processing at least one instruction (excluding semaphore waits). Units are in seconds. |
|
number (float) |
percent |
No |
Percentage of time when Sync engine is processing at least one instruction (excluding semaphore waits). |
|
integer (int64) |
No |
Number of Sync Engine instructions. |
|
|
number (float) |
seconds |
No |
Total duration of all GpSimd engine instructions. Overlapping instructions on the engine are added together. Units are in seconds. |
|
number (float) |
seconds |
No |
Duration of time when GpSimd engine is processing at least one instruction (excluding semaphore waits). Units are in seconds. |
|
number (float) |
percent |
No |
Percentage of time when GpSimd engine is processing at least one instruction (excluding semaphore waits). |
|
integer (int64) |
No |
Number of GpSimd Engine instructions. |
|
|
number (float) |
seconds |
No |
Total duration of all Collective Communication (CC) operations. Overlapping operations are added together. Units are in seconds. |
|
number (float) |
seconds |
No |
Duration of time when at least one Collective Communication (CC) operation is active. Units are in seconds. |
|
number (float) |
percent |
No |
Percentage of time when at least one Collective Communication (CC) operation is active. |
|
integer (int64) |
No |
Number of Collective Communication (CC) operations. |
|
|
number (float) |
seconds |
No |
Total duration of instructions on all Collective Communication (CC) cores. Overlapping instructions on the cores are added together. |
|
number (float) |
seconds |
No |
Duration of time when at least one Collective Communication (CC) core is processing an instruction (excluding semaphore waits). |
|
number (float) |
percent |
No |
Percentage of time when at least one Collective Communication (CC) core is processing an instruction (excluding semaphore waits). |
|
integer (int64) |
No |
Number of instructions on all Collective Communication (CC) cores. |
|
|
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. |
|
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. |
|
number (float) |
percent |
No |
Percentage of total execution time when at least one engine, DMA, or CC operation is active. |
SystemProfileEvents#
This table contains system profile events, which are generated from the Neuron Runtime, the framework code, or the Python application.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int64) |
Yes |
Unique identifier for this event. Stored as int64 in Parquet. The API returns this as a string via SQL CAST to prevent JavaScript Number precision loss for values > 2^53. generate_schema.sh post-processes the TypeScript output to emit string instead of number for this field. |
|
|
string |
Yes |
Name of the event. Example: |
|
|
integer (int64) |
nanoseconds |
Yes |
Timestamp when event started. |
|
integer (int64) |
nanoseconds |
Yes |
Timestamp when event finished. |
|
integer (int64) |
nanoseconds |
Yes |
Duration of this event. |
|
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: |
|
|
string |
Yes |
Unique identifier of the server instance that produced this event. Example: |
|
|
integer (int64) |
Yes |
The process id on that produced this event. Not necessarily unique when multiple instances or containers are in use. |
|
|
integer (int64) |
Yes |
The thread id on that produced this event. Not necessarily unique when multiple instances or containers are in use. |
|
|
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. |
|
|
integer (int64) |
No |
Index of the physical NeuronCore within a NeuronDevice. |
|
|
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. |
|
|
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: |
|
|
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: |
|
|
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: |
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.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int64) |
nanoseconds |
Yes |
Timestamp for this HBM memory utilization measurement. |
|
integer (int64) |
bytes |
Yes |
Total HBM bytes used across all usage types. |
|
integer (int64) |
bytes |
Yes |
HBM bytes used for Generic allocations. |
|
integer (int64) |
bytes |
Yes |
HBM bytes used for Tensor Instructions. |
|
integer (int64) |
bytes |
Yes |
HBM bytes used for Scalar Instructions. |
|
integer (int64) |
bytes |
Yes |
HBM bytes used for GpSimd Instructions. |
|
integer (int64) |
bytes |
Yes |
HBM bytes used for IO. |
|
integer (int64) |
bytes |
Yes |
HBM bytes used for DRAM Spill. |
|
integer (int64) |
bytes |
Yes |
HBM bytes used for Weights. |
|
integer (int64) |
bytes |
Yes |
HBM bytes used for Profiler Buffers. |
|
integer (int64) |
bytes |
Yes |
HBM bytes used for Scalar Table. |
|
integer (int64) |
bytes |
Yes |
HBM bytes used for Vector Instructions. |
|
integer (int64) |
bytes |
Yes |
HBM bytes used for Sync Instructions. |
|
integer (int64) |
bytes |
Yes |
HBM bytes used for Shared Scratchpad. |
|
integer (int64) |
bytes |
Yes |
HBM bytes used for Tensors. |
|
integer (int64) |
bytes |
Yes |
HBM bytes used for Ucode lib. |
|
integer (int64) |
bytes |
Yes |
HBM bytes used for GpSimd STDIO. |
|
integer (int64) |
bytes |
Yes |
HBM bytes used for Collectives. |
|
integer (int64) |
bytes |
Yes |
HBM bytes used for Scratchpad. |
|
integer (int64) |
bytes |
Yes |
HBM bytes used for DMA Rings Runtime. |
|
integer (int64) |
bytes |
Yes |
HBM bytes used for DMA Rings Spill. |
|
integer (int64) |
bytes |
Yes |
HBM bytes used for DMA Rings IO. |
|
integer (int64) |
bytes |
Yes |
HBM bytes used for DMA Rings Collectives. |
|
integer (int64) |
bytes |
Yes |
HBM bytes used for unrecognized usage types. |
|
integer (int64) |
Yes |
Neuron device index relative to the instance |
|
|
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. |
|
|
string |
Yes |
Identifier for the instance on which this HBM memory utilization was measured. Example: |
SystemProfileMetadata#
Metadata for a system profile.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
string |
Yes |
First timestamp in the system profile trace, in human-readable format. Example: |
|
|
string |
Yes |
Last timestamp in the system profile trace, in human-readable format. Example: |
|
|
string |
Yes |
First timestamp in the system profile trace, in Unix nanoseconds format. Example: |
|
|
string |
Yes |
Last timestamp in the system profile trace, in Unix nanoseconds format. Example: |
|
|
integer (int64) |
Yes |
Version of the NTFF (Neuron Trace File Format) used for this profile. Example: |
|
|
integer (int64) |
bytes |
Yes |
Total HBM capacity in bytes for a single NeuronCore. |
TensorInfo#
Information about tensors used in the NEFF.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
string |
No |
The name of the tensor. Example: |
|
|
string |
No |
What the tensor is used for. For example input tensor, output tensor, or weight tensor. Example: |
|
|
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: |
|
|
string |
No |
The multi-dimensional shape of the tensor. Example: |
|
|
integer (int64) |
bytes |
No |
The total size of the tensor in bytes. |
|
integer (int64) |
No |
Neff node. Example: |
|
|
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. |
|
|
integer (int64) |
No |
The average size in bytes of each DMA transfer when loading this tensor into the State Buffer. |
|
|
integer (int64) |
No |
The total size in bytes of all DMA transfers when loading this tensor into the State Buffer. |
|
|
integer (int64) |
No |
The total number of DMAs that loaded this tensor into the State Buffer. |
|
|
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. |
Throttle#
Hardware throttling notification events that track thermal or power throttling on NeuronCores.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int64) |
nanoseconds |
Yes |
Duration of the throttling event. |
|
integer (int32) |
Yes |
Allowed utilization level value used to calculate throttling limit (utilization_limit = k/n, throttling = 1 - utilization_limit). |
|
|
integer (int32) |
Yes |
Maximum possible utilization level value used to calculate throttling limit (utilization_limit = k/n, throttling = 1 - utilization_limit). |
|
|
integer (int64) |
nanoseconds |
Yes |
Timestamp when the throttling event occurred. |
|
string |
Yes |
Type of throttling (e.g., “activity_0”, “activity_1”, “unknown”). Example: |
|
|
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. |
ThrottleSummary#
Summary statistics for hardware throttling on each NeuronCore and for each type of throttling.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
integer (int32) |
Yes |
Index of the NeuronCore within a Neuron Device. |
|
|
string |
No |
Type of throttling (e.g., “activity_0”, “activity_1”, “unknown”). The value “all” means the metric aggregates across all throttle types. Example: |
|
|
integer (int64) |
nanoseconds |
No |
Total time in nanoseconds when throttling was active. |
|
number (float) |
percent |
No |
Percent of total time when throttling was active. |
|
number (float) |
percent |
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 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. |
Warning#
Warnings produced during profile processing. Used to communicate the possibility of missing or unexpected data.
Field |
Type |
Unit |
Required |
Description |
|---|---|---|---|---|
|
string |
Yes |
The message displayed to users to describe the warning. Example: |
|
|
string |
Yes |
The type of warning captured during profile processing. Example: |
Enums#
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.
Allowed values:
instructioninputoutputdatasoftware_dynamichardware_dynamicunknown
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 NotificationErrorsWarnings
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_devicedevice_to_hostbidirectional
PerformanceMode#
Performance mode for a Vector or Scalar Engine instruction. Indicates the throughput optimization applied by the hardware.
Allowed values:
2x4x
This document is relevant for: Inf1, Inf2, Trn1, Trn2, Trn3