NVIDIA Tools Extension Library (NVTX)

NVTX (NVIDIA Tools Extension Library) is a cross-platform API used to annotate source code with contextual information. This information is consumed by developer tools (like NVIDIA Nsight Systems) to provide meaningful insights during profiling.

Key Behavior: By default, NVTX API calls do nothing. They only become active when the program is launched from a developer tool that implements the NVTX API. When active, these calls can be used to:

• Record traces on a timeline.
• Build statistical profiles of time spent in specific code sections.
• Print messages to the console.
• Enable/disable specific tool features based on the current execution context.

NVTX annotations are natively supported by several NVIDIA profiling and analysis tools:

• Nsight Systems: Logs NVTX calls and displays them on a timeline alongside driver, OS, and hardware events.
• Nsight Compute: Uses NVTX ranges to focus deep-dive GPU performance analysis on specific code sections.
• Nsight Graphics: Uses NVTX ranges to set bounds for range profiling within the Frame Debugger.
• CUPTI API: Supports recording traces of NVTX calls.

NVTXW ("NVTX Writer") is an offline C API used to deliver pre-collected trace data (events, counters, and timing information) to tools via an NVTXW backend.

Unlike standard NVTX, which performs online instrumentation by intercepting calls during a live process, NVTXW is designed for programmatic offline import.

Use NVTXW when:

• You have custom profiler/framework data that you want to render on a tool's timeline (e.g., Nsight Systems).
• You need to replay a previously captured trace.
• You are converting third-party trace files (like OTF2) into NVTX reports.
• You want to merge foreign trace data into an existing profiler report.
• You are building a synthetic-report generator for testing or demos.

In the NVTXW model, you supply timestamps explicitly rather than relying on the tool to capture them at the moment of the API call.

NVTXW provides specific thread-safety guarantees that producers must follow:

• Concurrent Writes: Distinct streams can be written to concurrently from different threads without requiring producer-side synchronization.
• Single Stream Writes: You should not assume that concurrent writes to a single stream are safe.
• Lifecycle and Registration: Calls related to lifecycle and registration (e.g., SessionBegin, StreamOpen, StreamClose, SessionEnd, and all *Register functions) are not assumed to be thread-safe. These must be serialized and should not overlap with data writes.

NVTX is designed for the following environments:

• Operating Systems: Windows, Linux, and other POSIX-like platforms (including cygwin), and Android.
• Process Bit-width: Both 64-bit and 32-bit processes are supported.
• CPU Architecture: No restrictions.

Important Limitations:

• No GPU Support: NVTX is not supported in GPU code (e.g., __device__ functions in CUDA). For efficient instrumentation of CUDA GPU code, use the pmevent PTX instruction.
• Dynamic Library Requirement: NVTX requires the platform's standard API to load a dynamic library (.dll or .so).

Overhead Behavior

• When no tool is present: The first NVTX call initializes the library and disables all functions. Subsequent calls become lightweight, inlined instructions that jump over the disabled call.
• When a tool is present: Initialization configures the API to jump directly into the tool's implementation. The overhead is then determined by the tool itself.

Managing Initialization Latency

The first NVTX call can incur significant overhead while loading the tool. If your program is latency-sensitive (e.g., a game), this first call might cause unexpected behavior. To avoid this, use the nvtxInitialize C API function to force-initialize NVTX at a convenient time during program startup.

To prevent annotation data from different libraries from clashing, each library should create and use its own dedicated nvtx.Domain. This allows profiling tools to group data by library and enables users to enable or disable specific library annotations during execution.

Use hierarchical category names (e.g., filesystem/path/to/category) within a domain to organize annotations into a logical tree structure that tools can navigate.

The NVTX parser uses a Streaming Visitor model via the PayloadStreamVisitor interface. Instead of building an intermediate data structure, it emits a depth-first stream of typed events. This ensures constant memory usage regardless of payload complexity and allows for incremental output generation.

To implement a custom processor, you can implement the PayloadStreamVisitor interface and handle the following event sequence:

OnBeginPayloads(count)
  OnPayloadBegin(index, schemaId, size, name)
    OnFieldBegin(name, description)
      [OnArrayBegin(length)]
        OnSignedInteger / OnUnsignedInteger / OnFloatingPoint / 
        OnString / OnRawBytes / 
        OnObjectBegin ... OnObjectEnd
      [OnArrayEnd]
    OnFieldEnd
    ...
  OnPayloadEnd
  ...
OnEndPayloads

The NvtxPayloadParser::ProcessPayloads method is the entry point for decoding NVTX payloads. It processes nvtxPayloadData_t elements by:

1. Validating the payload pointer, size, and schema ID.
2. Resolving SIZE_MAX payloads for null-terminated C-string types by scanning for the terminator.
3. Routing predefined types (where schema ID is < NVTX_PAYLOAD_SCHEMA_ID_STATIC_START) to EmitPredefinedTypePayload. If the payload size is a multiple of the element size, it is automatically detected as a fixed-size array.
4. Looking up registered and finalized schemas for custom IDs, then iterating through fields via VisitSchemaFields.