NTFF Trace File Format (ntff.proto)

All addresses, offsets, and field numbers on this page apply to libnrt.so from aws-neuronx-runtime-lib 2.31.24.0-0b044f4ce (build-id 8bb57aba0fb2e0035f1d88e9fc4fb3e7387c102e, package KaenaProfilerFormat-2.31.0.0). The ELF is not stripped; full .symtab + DWARF are present. .text/.rodata VMA equals file offset (.rodata base 0x7cf000), and .data here is also VMA==file-offset (0xc07e00), so every xxd -s 0x… on a .rodata/.data-resident descriptor reads the cited byte. Other versions will differ. Evidence grade: Confirmed (byte-anchored) — the entire schema is recovered from the two embedded FileDescriptorProto blobs in .rodata, decoded with the protobuf descriptor parser; every message, field number, type, label, oneof, and enum below is the protoc-input definition, not an inference. · Part XIII — Profiling, Trace & Telemetry · back to index

Abstract

A NeuronCore profile is serialized as NTFF — the Neuron Trace File Format, an AWS-authored protobuf schema generated by protoc against vendored google::protobuf 26.1 (GOOGLE_PROTOBUF_VERSION 5026001) + Abseil LTS 20230802, and compiled directly into libnrt.so. The schema lives in two .proto files in one package "ntff": ntff.proto (the device-profile schema rooted at ntff_info) and neuron_trace.proto (the host system-trace schema rooted at ntrace_info, which import "ntff.proto"). Together they declare 40 message classes, 9 + 3 file-level enums, and 207 fields (205 scalar/message fields + 2 map<> fields). This is the authoritative on-disk contract for everything the profiler emits: the per-NC .ntff device-profile files, ntrace.pb, cpu_util.pb, and host_mem.pb.

The recovery technique is what makes this page ground truth rather than reconstruction. Each .proto file is compiled into a serialized FileDescriptorProto blob stored verbatim in .rodata — descriptor_table_protodef_ntff_2eproto @0xad4200 (7185 B) and descriptor_table_protodef_neuron_5ftrace_2eproto @0xad6760 (2570 B). Rather than inferring the wire schema from ByteSizeLong offsets, the descriptor parser decodes those blobs back into the exact source schema. Every field number, name, type, label, oneof, and enum value below is therefore the protoc input, byte-for-byte. The pointers and sizes come from the two TcParseTableBase descriptor tables in .data (descriptor_table_ntff_2eproto @0xc0b6a0, …neuron_5ftrace… @0xc0c0c0); the 40-class roster matches the RTTI vtable band 0xbf6fa8..0xbf8708 exactly.

The familiar reference frame is any protoc-generated proto3 file pair with a cross-file import. NTFF's twists are three: the schema is split across two files that share one package (so neuron_trace.proto references ntff.proto types like neff_arch and ultraserver_info directly); subgraph_info is a 17-field hub message that every device-profile node tree funnels through; and the container is not a bare protobuf stream — a per-NC .ntff file carries a 128-byte header (byte0 = format-version 6, bytes[1..8] = ByteSizeLong) followed by the ntff_info message and then raw instruction/notification blob sections referenced by byte offset from inside the proto.

For reimplementation, the contract is:

• The 40-message schema — all ntff:: messages, their field numbers/types/labels/oneofs, and the three enums that gate the device-profile (notification_trace_type, block_type, memory_usage_type) plus the three host-trace enums (ntrace_data_type, ntrace_event_phase_type, ntrace_message_type).
• The containment tree and three roots — ntff_info (device profile), ntrace_info (host trace), and the two standalone record streams cpu_util / host_mem_usage — and the three oneof groups (one of which is a proto3 synthetic-optional).
• The on-disk NTFF container framing — the 128-byte header on .ntff device-profile files, the SerializeToOstream body, and the trailing instruction/notification blob sections; versus ntrace.pb, which is a bare ntrace_info protobuf stream with no header.
• The notification_type → notification_trace_type + block_type map — the device-NQ converter that wraps each drained notification into an ntff::trace_info.

1. The Schema Recovery Model

Purpose

NTFF is reverse-engineered by recovering the source schema, not by reconstructing it from serializer logic. proto3 stores the entire schema — every message, field name, number, type, label, oneof grouping, and enum — inside a single serialized FileDescriptorProto, which protoc embeds verbatim in the generated *.pb.cc translation unit. libnrt.so contains both blobs in .rodata, so the schema is recoverable to byte accuracy by reading those bytes and handing them to the descriptor parser. This page's §2/§3 are the decoded output of exactly that operation.

NOTE — the alternative — inferring field numbers and wire types from ByteSizeLong / _InternalSerialize tag bytes — is what the sibling NTFF Wire Tables page does to cross-validate this schema (zero contradictions across all 38 message tables). The schema here is the primary source; the wire tables are the independent check.

Entry Point

The descriptor tables are the on-disk handles that point at the two blobs. Each is a google::protobuf::internal::DescriptorTable in .data:

descriptor_table_ntff_2eproto (0xc0b6a0, .data)            ── ntff.proto handle
  +0x04 serialized_size = 0x1c11 (7185)
  +0x08 char* protodef  = 0xad4200  ── descriptor_table_protodef_ntff_2eproto (.rodata)
  +0x10 const char* filename = 0x8487b7 -> "ntff.proto"
  +0x28 int num_messages = 0x1f (31)
  +0x48 file_level_metadata = 0xcafe00 (.bss, Metadata[31], lazy AssignDescriptors)

descriptor_table_neuron_5ftrace_2eproto (0xc0c0c0, .data)  ── neuron_trace.proto handle
  +0x04 serialized_size = 0x0a0a (2570)
  +0x08 char* protodef  = 0xad6760  ── descriptor_table_protodef_neuron_5ftrace_2eproto
  +0x10 filename = 0x848866 -> "neuron_trace.proto"
  +0x20 deps      = 0xbf8860 -> deps[0] = 0xc0b6a0 (ntff.proto)   ── 1 dependency
  +0x28 num_messages = 0x09 (9)  ── 7 top-level + 2 map-entry

Algorithm

The recovery is a four-line operation; the only subtlety is that .rodata VMA equals file offset here, so the descriptor-table pointer doubles as a file offset.

// Recover the NTFF schema from the embedded FileDescriptorProto blobs.
function recover_ntff_schema(elf):
    dt_ntff  = read_struct(elf, 0xc0b6a0)          // descriptor_table_ntff_2eproto (.data)
    blob     = read_bytes(elf, dt_ntff.protodef,   // 0xad4200, len 0x1c11 (=7185)
                                dt_ntff.serialized_size)
    fdp_ntff = FileDescriptorProto.parse(blob)     // -> 31 msgs, 9 enums, 155 fields
    dt_tr    = read_struct(elf, 0xc0c0c0)          // descriptor_table_neuron_5ftrace_2eproto
    blob2    = read_bytes(elf, dt_tr.protodef,     // 0xad6760, len 0x0a0a (=2570)
                                dt_tr.serialized_size)
    fdp_tr   = FileDescriptorProto.parse(blob2)    // -> 7 top msgs (+2 map-entry), 3 enums, 52 fields
    // declaration order in each FDP == file_level_metadata index order == source order.
    return (fdp_ntff, fdp_tr)

The blob signatures are self-evident in the first bytes (both byte-verified at the cited offsets):

0xad4200: 0a 0a "ntff.proto" 12 04 "ntff" 22 88 01 0a 18 "notification_buffer_size" …
          └field1 name──────┘ └field2 pkg─┘ └field4 message, len 0x88, name 24B┘
0xad6760: 0a 12 "neuron_trace.proto" 12 04 "ntff" 1a 0a "ntff.proto" 22 c3 03 0a 0c "ntrace_event" …
          └field1 name───────────────┘ └field2 pkg─┘ └field3 dependency┘ └field4 msg, name 12B┘

QUIRK — field declaration order inside a FileDescriptorProto is source order, not field-number order. A handful of NTFF messages declare fields out of numeric order on purpose — engine_instruction_info declares block_type=5 and function_section=6 before nc_engine_type=1; neff_node_info puts uid=5 after its node_info oneof; trace_info and ntrace_event similarly. The field numbers below are authoritative for the wire; the declaration order is preserved only as a faithful echo of the source.

2. ntff.proto — the Device-Profile Schema

Purpose

ntff.proto (package ntff, 31 messages, 9 enums, 155 fields) is the schema of a per-NeuronCore device-profile .ntff file. Its root is ntff_info. The device-profile producer drains the on-device Notification Queues, converts each notification into a trace_info, populates the subgraph_info node tree, and serializes one ntff_info per NC. The collectives sub-tree (collectives_*, 14 messages) mirrors the live nccom DMA path.

Enums

Nine file-level enums gate the device-profile messages. The two with non-obvious numbering are flagged.

memory_usage_type (verbatim, 18 values): 0 GENERIC, 1 INSTR_PE, 2 INSTR_ACT, 3 INSTR_POOL, 4 IO, 5 DRAM_SPILL, 6 WEIGHT, 7 DMA_RING, 8 NOTIFICATION, 9 ACT_TABLE, 10 INSTR_DVE, 11 INSTR_SP, 12 ONE_TMPBUF, 13 TENSOR, 14 UCODE_LIB, 15 POOL_STDIO, 16 CC, 17 ONE_TMPBUF_NOT_SHARED.

GOTCHA — dma_queue_type skips value 5. A reimplementation that derives the enum from a dense 0..N loop will mis-assign H2T to 5 instead of 6 and silently corrupt every dma_queue_usage_info.queue_type. The descriptor stores H2T = 6 explicitly.

The root and node tree

The device-profile message set is dominated by the ntff_info root and the subgraph_info hub. Both are large; they are presented field-by-field because every field is a distinct wire contract.

ntff_info (descriptor idx 1, 20 fields — the .ntff root):

subgraph_info (descriptor idx 19, 17 fields — the node-tree hub, filled by _populate_subgraph_info):

The remaining 29 messages — by group

The other ntff.proto messages are small (1–9 fields each). Presented as grouped dimension tables; field numbers run 1..N in declaration order unless a number is called out, and the Note column flags every non-dense numbering, oneof, or packed repeated.

Profile metadata + timeline (8 messages):

Per-engine instruction + trace (6 messages):

Instruction patch + DMA usage (3 messages):

Collectives telemetry (12 messages — mirrors the live nccom DMA path):

GOTCHA — collectives_channel has no field 2 (name=1, subchannels=3, engine_index=4). Like dma_queue_type's missing value 5, this is an intentional schema gap from a removed field; the wire is correct, a dense reimplementation is not.

3. neuron_trace.proto — the Host-Trace Schema

Purpose

neuron_trace.proto (package ntff, import "ntff.proto", 7 top-level messages + 2 auto-generated map<> entries, 3 enums, 52 fields) is the host system-trace schema. Its root is ntrace_info (serialized to ntrace.pb); it also defines two standalone record streams cpu_util and host_mem_usage (written to cpu_util.pb / host_mem.pb). Because it imports ntff.proto, it references ntff::neff_arch and ntff::ultraserver_info directly across the file boundary — wired in-binary by descriptor_table_neuron_5ftrace_2eproto_deps[0] @0xbf8860 = 0xc0b6a0 (the ntff.proto table).

Enums

NOTE — ntrace_data_type stops at 14 in the proto; the C-side ntrace_data_type enum (the in-process ntrace_data_t ring, event-taxonomy) adds 15 = COUNT as a sentinel. The sentinel is a producer-side bound, never serialized, so it is correctly absent from the wire schema.

Messages

ntrace_info and ntrace_event are large; the rest are small. The two map fields produce the auto-generated AttributesEntry/IdMapEntry map-entry messages (the 39th and 40th classes).

ntrace_info (descriptor idx 4, 18 fields — the ntrace.pb root):

ntrace_event (descriptor idx 0, 16 fields — one host trace event):

The remaining 5 top-level + 2 map-entry messages:

The two map-entry signatures are byte-confirmed from the RTTI typeinfo names: MapEntry<…AttributesEntry…, string, string, FieldType 9, FieldType 9> and MapEntry<…IdMapEntry…, unsigned long, string, FieldType 4, FieldType 9> (FieldType 4 = uint64, 9 = string).

Containment, roots, and oneofs

Four messages are roots — serialized as file-level objects, never embedded:

ntff_info        ── per-NC .ntff device profile (128-byte header + this; §4)
ntrace_info      ── ntrace.pb host trace (bare protobuf stream, no header)
cpu_util         ── cpu_util.pb (standalone record stream)
host_mem_usage   ── host_mem.pb (standalone record stream)

The device-profile tree funnels through subgraph_info; the host-trace tree through ntrace_info. There are exactly three oneof groups across both files:

QUIRK — the only proto3 optional field in the entire schema is ntrace_event.nrta_seq_id (#16). proto3 implements explicit presence by wrapping the field in a one-member synthetic oneof named _nrta_seq_id with proto3_optional=true in the descriptor. A reimplementer reading the raw FileDescriptorProto sees a third oneof; it is a presence marker, not a variant selector — on the wire it is a plain uint64 whose has-bit is tracked. Every other field is proto3 implicit-presence (no has-bit; default-valued fields are not serialized).

4. The On-Disk NTFF Container

Purpose

NTFF defines two distinct on-disk containers. The per-NC device-profile .ntff file is not a bare protobuf stream: it carries a 128-byte fixed header, then the ntff_info message, then raw blob sections that the proto references by byte offset. The host ntrace.pb (and cpu_util.pb / host_mem.pb) are bare protobuf streams written straight from SerializeToOstream with no framing. Confusing the two is the single most common container-level error.

Entry Point

nrt_inspect_stop (0x9f130)                              ── fused harvest + serialize
  ├─ per-NC device profile:
  │    nrt_profile_serialize (0xb2670)                  ── .ntff: 128-B header + ntff_info + blobs
  │    nrt_profile_session_serialize (0xb61d0)          ── session variant (delegates to the above)
  └─ host trace:
       nrt_trace_serialize (0xae300)                    ── ntrace.pb: BARE ntrace_info stream
       datafile_init (0x9ac30) x3                        ── cpu_util.pb / host_mem.pb / ntrace.pb descriptors

Algorithm — the .ntff device-profile framing

nrt_profile_serialize @0xb2670 writes the header from a 128-byte stack buffer, then the message, then two blob loops. The header layout and version byte are byte-confirmed in the decompile (char v46[128], v46[0] = 6, *(_QWORD*)&v46[1] = ByteSizeLong).

// Models nrt_profile_serialize @0xb2670.
// args: path, ntff_info* info, vector<{char*,u32}>* instruction_data, *trace_notifications
function nrt_profile_serialize(path, info, instr_blobs, trace_blobs):
    if !is_valid_directory(dirname(path)): return FAIL    // nlog "...not a valid directory"
    mkdir_recursive(dirname(path))                         // utils_mkdir_recursive

    char header[128] = {0}                                 // v46[128], zero-filled
    header[0]    = 6                                       // FORMAT VERSION byte  (v46[0] = 6)
    *(u64*)&header[1] = info->ByteSizeLong()               // serialized proto length, offset +1
    // header[9..127] reserved (zero)

    ofs = ofstream(path)
    ofs.write(header, 128)                                 // 128-byte header
    if !info->SerializeToOstream(&ofs):                    // the ntff_info protobuf body
        nlog("Failed to serialize a profile file at: %s ...")
        return FAIL
    for (ptr,len) in *instr_blobs:  ofs.write(ptr, len)    // raw instruction blob section
    for (ptr,len) in *trace_blobs:  ofs.write(ptr, len)    // raw notification blob section
    return SUCCESS

The file therefore has four regions:

+--------------------------------------------------------------+  offset 0
| HEADER (128 bytes)                                           |
|   [0]     u8   format_version = 6                            |
|   [1..8]  u64  proto_byte_size  (= ntff_info.ByteSizeLong()) |
|   [9..127]     reserved (zero)                               |
+--------------------------------------------------------------+  offset 128
| ntff_info  (serialized protobuf, length = proto_byte_size)   |
+--------------------------------------------------------------+  offset 128 + proto_byte_size
| instruction blob section  (concatenated raw byte runs)       |  <- engine_instruction_info
|                                                              |     .instructions_offset / _size
+--------------------------------------------------------------+
| notification blob section (concatenated raw byte runs)       |  <- trace_info.trace_offset / _size
+--------------------------------------------------------------+  EOF

QUIRK — the bulky payloads — engine instruction streams and drained NQ notification blobs — are not inlined into the proto as bytes fields. The proto carries only descriptors: engine_instruction_info.instructions_offset/_size and trace_info.trace_offset/_size are byte offsets into the trailing blob sections, relative to the start of each section. This keeps the protobuf body small and lets the consumer mmap the raw blobs. A reader that expects the instruction bytes inside the instructions field will find it empty when the offset/size descriptors are populated instead — the two are alternative encodings of the same payload.

The bare ntrace.pb stream

nrt_trace_serialize @0xae300 writes no header: it opens the file and calls SerializeToOstream(ntrace_info) directly. The decompile shows only the SerializeToOstream call guarded by the "Failed to serialize a trace file at: %s" nlog — there is no [0]=6 header byte and no write(buf,128). cpu_util.pb and host_mem.pb are likewise bare record streams; their existence and num_events are advertised through ntrace_data_file descriptors inside ntrace_info / host_stats rather than through any file header.

5. The 205-vs-207 Field Tension

The schema declares a field count that depends on how map<> fields are counted, and the two NTFF wiki pages report the two different numbers on purpose. This section reconciles them so a reimplementer is not surprised.

The descriptor parse yields 207 top-level FieldDescriptorProto records (155 in ntff.proto + 52 in neuron_trace.proto). Two of those are the map<> fields — ntrace_event.attributes (#12) and interned_data_db.id_map (#2). A map<K,V> field is one FieldDescriptorProto (a repeated message field whose type is the auto-generated *Entry message), so it counts as one top-level field but expands into an extra message with its own key/value pair.

ntff.proto              : 31 messages, 155 fields, 0 map fields
neuron_trace.proto      :  7 messages,  52 fields, 2 map fields, 1 proto3-optional
                          ──────────────────────────────────────
top-level FieldDescriptorProto records              = 207
  minus the 2 map<> fields (counted as map decls)   = 205   <- this page's §2/§3 scalar/message count
map-entry messages (AttributesEntry, IdMapEntry)    =   2  (the 39th/40th classes)
  each contributing key+value FieldEntry records    = + 4

The sibling NTFF Wire Tables page byte-decodes the TcParseTableBase fast-parse statics and reports 207 FieldEntry records across the 38 top-level message tables, with zero type mismatches against this schema. That 207 is the same 205 scalar/message fields plus the 2 map fields, because in the parent table a map<> field is one FieldEntry (it dispatches to a FastMtR1/MapAuxInfo handler that owns the nested entry). The map-entry messages' own key/value (+4) live in their own two parse tables, which is why the wire-tables total is 207 (parent tables) and not 211.

CORRECTION (W2-NTFF) — an early reading conflated "205 fields" (the §2/§3 scalar/message declarations) with "207 FieldEntry records" (the wire-table decode) as a contradiction. They are not: 207 = 205 + the 2 map<> fields, which §2/§3 list as map declarations rather than as scalar fields. The two numbers describe the same schema at two granularities. The grand class count (40 = 31 + 7 + 2 map-entry) is invariant under either convention.

6. The notification_type → trace_info Map

Purpose

The device-profile producer drains the on-device Notification Queues; each raw notification carries a notification_type (0..10). Before it can be wrapped in an ntff::trace_info, that type is converted into an ntff::notification_trace_type (the trace_info.trace_type field) plus an ntff::block_type (the trace_info.block_type field). The converter is nrt_profile_convert_trace_type_from_ntff_params @0xaaf40, dispatched through a 9-entry jump table @0x8570c0.

Algorithm

// Models nrt_profile_convert_trace_type_from_ntff_params @0xaaf40.
// out: *trace_type (notification_trace_type), *block_type (block_type); ret NRT_STATUS.
function convert_trace_type(notification_type, &trace_type, &block_type):
    switch (notification_type):                         // jump table @0x8570c0, range 0..8
        case 0 TRACE:        trace_type=0 INSTRUCTION; block_type=0 NC;    return SUCCESS
        case 1 EVENT:        trace_type=1 EVENT;       block_type=0 NC;    return SUCCESS
        case 2 ERROR:        trace_type=2 ERROR;       block_type=0 NC;    return SUCCESS
        case 3 INFER_STATUS: trace_type=0 INSTRUCTION; block_type=0 NC;    return SUCCESS
        case 4 DMA:          trace_type=4 DMA;         block_type=1 DMA;   return SUCCESS
        case 5 THROTTLE:     trace_type=5 THROTTLE;    block_type=0 NC;    return SUCCESS
        case 6 TOPSP_TRACE:  trace_type=0 INSTRUCTION; block_type=2 TOPSP; return SUCCESS
        case 7 TOPSP_EVENT:  trace_type=1 EVENT;       block_type=2 TOPSP; return SUCCESS
        case 8 TOPSP_ERROR:  trace_type=2 ERROR;       block_type=2 TOPSP; return SUCCESS
        default:  // 9 TOPSP_CC_STATUS, 10 MAX
            nlog("Invalid notification type: %d");      // @0x83e9b7
            return NRT_INVALID                          // = 2; notification dropped

The full mapping, byte-verified against both the jump table and the decompiled switch (independent agreement):

NOTE — notification_trace_type = 3 EXPLICIT is never produced by this converter. It is reserved for the explicit collectives-status path (nrt_profile_session_append_cc_notifications @0xaf700), which writes trace_info.trace_type = EXPLICIT directly. So the converter's domain is notification_type 0..10 (only 0..8 mapped) and its codomain is notification_trace_type 0,1,2,4,5 — value 3 is filled by a different code line. A reimplementation that routes every trace_info through this converter will never emit EXPLICIT traces.

Each converted notification becomes one trace_info appended to ntff_info.session_notifications (or subgraph_info.traces), with the raw notification bytes pushed into the blob section and referenced by trace_info.trace_offset/_size.

Verification

Every claim on this page is anchored to the live binary (build-id 8bb57aba…). The schema is not inferred from serializer logic — it is the decoded FileDescriptorProto:

• Both FDP blobs decode with zero errors. ntff.proto @0xad4200 (7185 B) → 31 messages, 9 enums, 155 fields; neuron_trace.proto @0xad6760 (2570 B) → 7 top messages, 3 enums, 52 fields (incl. 2 map<>, 1 proto3-optional). The descriptor-table pointers/sizes are byte-read from .data: descriptor_table_ntff_2eproto @0xc0b6a0 has +0x04 = 0x00001c11 and +0x08 = 0x00000000_00ad4200; the dependency …neuron_5ftrace…deps[0] @0xbf8860 = 0xc0b6a0.
• 40 classes match the RTTI roster. nm | rg _ZTVN4ntff returns exactly 40 vtables in the contiguous band 0xbf6fa8..0xbf8708 = 31 (ntff) + 7 (neuron_trace top) + 2 (AttributesEntry/IdMapEntry). The two map-entry typeinfo names confirm key/value FieldType 9,9 and 4,9.
• The container framing is byte-confirmed in the decompile. nrt_profile_serialize @0xb2670: char v46[128], v46[0] = 6, *(u64*)&v46[1] = ntff_info::ByteSizeLong(), ostream::write(…, v46, 128), then SerializeToOstream, then the two blob loops. nrt_trace_serialize @0xae300 has no header write — bare SerializeToOstream(ntrace_info).
• The NQ converter is dual-sourced. Jump table @0x8570c0 (range 0..8) and decompiled switch @0xaaf40 agree on all 9 mapped cases and the NRT_INVALID default for 9/10.
• Independent cross-validation. The NTFF Wire Tables TcParseTableBase decode reproduces all 207 parent-table FieldEntry records with zero type mismatches against this schema, and rebuilds the entire containment graph from the parse-table aux pointers alone.

Related Components

Cross-References

• NTFF Wire Tables (TcParseTable Decode) — the TcParseTableBase byte-decode that independently confirms all 207 field entries, wire types, and the containment graph; owner of the 207-vs-205 wire-side count
• SysTraceEventType Taxonomy (46 Variants) — the 46-variant sys_trace event model that feeds ntrace_event via event_to_proto, and the C ntrace_data_type ring with its 15 = COUNT sentinel
• The Inspect / Profile API — the nrt_inspect lifecycle that arms the producers and calls the container writers
• Overview: the Three Trace Producers — how device-profile, sys_trace, and the ntrace ring all serialize into this one schema family
• C++ Class Hierarchy and RTTI — the _ZTVN4ntff vtable band that pins the 40-class roster and the si:google::protobuf::Message base
• back to index