Instrumentation Profiler

Instrumentation profiling provides exact, deterministic measurements of function calls, loop iterations, and memory accesses by inserting timers or counters at specific code points. Unlike statistical sampling, it captures every instance of the instrumented event, enabling precise analysis of execution paths, call counts, and cumulative time. This is critical for understanding the fine-grained behavior of NPU kernels, such as the exact cost of a memory load within a tight loop or the overhead of a specific API call.

The primary trade-off of instrumentation is runtime overhead. The inserted measurement code (probes) adds instructions that would not otherwise be executed, slowing down the target program. This probe effect can distort performance measurements, especially for fine-grained operations. The overhead is proportional to the frequency of instrumented events. For NPU profiling, careful placement is required to avoid significantly altering kernel execution timing, which could invalidate optimization decisions based on the profiled data.

• Compile-Time Instrumentation: Probes are inserted by the compiler during the build process. This offers deep integration but requires source code and recompilation. Tools like gcc -pg use this method.
• Runtime Instrumentation (Binary Rewriting): Tools like DynamoRIO or Intel Pin modify the program's binary or intermediate representation in memory during execution. This allows profiling of closed-source binaries but can add higher overhead.
• Source-Level Instrumentation: Manual or preprocessor-driven insertion of profiling macros (e.g., LOG_TIME) into the source code offers maximum control but is labor-intensive.

Instrumentation and sampling profilers are complementary. Sampling periodically interrupts the program (e.g., via OS timers) to record the call stack, providing a low-overhead, statistical view of where time is spent. It identifies hotspots but can miss brief, frequent functions. Instrumentation provides complete event traces and exact counts but with higher overhead. A robust NPU performance analysis workflow often uses sampling to find general bottlenecks, then employs targeted instrumentation for deep, precise analysis of critical code sections.

Instrumentation typically generates two key data artifacts:

• Execution Traces: Chronological, event-by-event logs (e.g., function entry/exit, memory allocations). These are voluminous but enable detailed replay and analysis of complex interaction sequences.
• Call Graphs: Aggregated representations of function call relationships and cumulative time spent in each function and its descendants. This helps visualize the program's execution hierarchy and identify costly call paths. Tools like gprof generate flat profiles and call graphs from instrumented binaries.

Within NPU acceleration, instrumentation profiling is a core component of auto-tuning systems. It provides the ground-truth performance data (e.g., kernel runtime, cache miss rates) needed to evaluate different kernel configurations (tile sizes, unroll factors). While too heavyweight for evaluating every candidate in a large search space, it is often used to profile a representative subset to train a performance model. This model then predicts outcomes for other configurations, guiding the search efficiently. The precise data from instrumentation is crucial for building accurate models.

Tools that insert profiling hooks during the compilation of a computational graph or kernel for an NPU. This is a source-level transformation where the compiler adds calls to timing functions or performance counter APIs.

• Key Mechanism: Modifies the Intermediate Representation (IR) or final assembly to inject measurement code.
• Example: The LLVM compiler infrastructure can be used to insert calls to a custom profiling runtime before and after targeted loops or functions.
• Advantage: Provides precise control over what is measured, enabling function-level or even basic-block-level granularity.
• Use Case: Essential for building custom, low-overhead profilers integrated into a vendor's SDK or a proprietary compilation pipeline.

Pre-built libraries that a program links against to automatically instrument API calls, kernel launches, and memory operations. These libraries intercept calls to the NPU driver or runtime (e.g., OpenCL, Vulkan, vendor-specific APIs).

• Key Mechanism: Uses function wrapping or dynamic linking to trap calls to functions like clEnqueueNDRangeKernel.
• Examples: NVIDIA's NVTX (NVIDIA Tools Extension) for annotating code regions, or custom wrappers built with LD_PRELOAD on Linux.
• Output: Generates timelines showing kernel execution, memory transfers, and host-device synchronization.
• Primary Benefit: Requires no source code modification, making it ideal for profiling closed-source libraries or full application stacks.

The foundational layer of instrumentation profiling, providing direct access to the NPU's internal performance monitoring units (PMUs). These counters track low-level hardware events.

• Measured Events: Includes compute throughput (FLOPS, TOPS), memory bandwidth utilization, cache hit/miss rates, pipeline stalls, and thread divergence.
• Tool Role: Profiling frameworks provide APIs or wrappers to configure, start, stop, and read these counters for specific kernel executions.
• Critical Insight: Reveals whether a kernel is compute-bound or memory-bound, guiding optimization efforts.
• Vendor Dependency: Access is provided through vendor-specific SDKs, such as ARM's Streamline Performance Analyzer for Ethos NPUs or Intel VTune for GPUs/NPUs.

Instrumentation profilers typically operate in one of two fundamental modes, each with a trade-off between detail and overhead.

• Tracing (Event-Based): Records a timestamped log of every instrumented event (e.g., function entry/exit). This creates a precise execution trace but can generate large amounts of data and add significant overhead.
• Statistical (Sampling): Periodically interrupts the program (or samples the program counter) to record the current call stack. This has lower overhead and identifies hotspots but can miss short, infrequent functions.
• Hybrid Approach: Advanced tools use instrumentation for key regions and sampling for broader context, balancing detail with system impact.

Instrumentation profilers do not operate in isolation; they are the data source for auto-tuning feedback loops. The profiler measures the outcome of each kernel configuration tried during a parameter search.

• Workflow: 1. A kernel tuner generates a kernel variant with a specific workgroup size or tile size. 2. The instrumented kernel is executed. 3. The profiler collects metrics (execution time, counters). 4. Results feed a performance model or search algorithm (e.g., Bayesian Optimization) to select the next configuration.
• Automation: This closed-loop process automatically explores the configuration space to find the optimal parameters for a given NPU and input size.

Raw profiling data is voluminous and complex. Dedicated visualizers transform this data into actionable insights, such as Flame Graphs for CPU/host time and timeline views for NPU activity.

• Timeline Views: Show concurrent kernel execution, memory transfers, and pipeline stalls, helping identify resource contention and idle periods.
• Flame Graphs: Visualize hierarchical call stacks, quickly pinpointing the hotspot identification within the host code driving the NPU.
• Metric Correlation: Cross-reference performance counter data (e.g., low cache hit rate) with specific lines of source code.
• Examples: Perfetto for system-wide tracing, NVIDIA Nsight Systems, AMD ROCprofiler, and ARM Mobile Studio.

A profiling method that periodically interrupts program execution (samples) to record the current call stack and program counter. It provides a statistical view of performance with very low overhead, making it suitable for production environments.

• Key Contrast: Unlike an Instrumentation Profiler, which inserts precise measurement code, a sampling profiler approximates hotspots by statistical inference.
• Overhead: Typically <1-5%, compared to higher overhead from instrumentation.
• Use Case: Initial, high-level identification of performance bottlenecks in long-running NPU applications.

A specialized software tool that measures the execution time, resource utilization, and hardware performance counters of a computational kernel running on an NPU or GPU.

• Scope: Focuses on the low-level, parallelized function (kernel) executing on the accelerator hardware.
• Metrics: Measures compute throughput, occupancy, memory bandwidth utilization, and pipeline stalls.
• Integration: Often part of vendor SDKs (e.g., NVIDIA Nsight, AMD ROCprofiler) and provides the detailed data needed for auto-tuning kernel parameters.

A hardware register within a processor that counts low-level architectural and microarchitectural events. These are the raw data sources for detailed performance analysis.

• Examples: Cache misses, floating-point operations (FLOPs), branch mispredictions, DRAM accesses, tensor core operations.
• Profiler Role: An Instrumentation Profiler or Kernel Profiler reads these counters to attribute costs to specific code regions.
• Critical Use: Determining if a kernel is compute-bound (high FLOP count) or memory-bound (high cache miss count).

A chronological, detailed record of the sequence of instructions, memory accesses, function calls, and kernel launches executed by a program. It captures temporal relationships and dependencies.

• Detail Level: Extremely high; can generate massive data files. Often used for debugging complex concurrency issues.
• Analysis: Used to visualize pipeline stalls, thread divergence, and asynchronous execution overlaps.
• Tool Relationship: An Instrumentation Profiler can be configured to generate a trace, but tracing is more exhaustive and has higher overhead than standard profiling.

A visualization of hierarchical profiler data, where the width of horizontal bars represents the proportion of time or samples spent in each function call stack.

• Purpose: To rapidly identify performance hotspots across the entire call hierarchy in a single view.
• Data Source: Can be generated from both sampling profiler data and instrumentation profiler data.
• Key Feature: Allows quick navigation from high-level application functions down to specific, costly NPU kernel calls.

The systematic process of using profiling data to identify the primary limiting factor (bottleneck) that restricts the overall performance of a system or application.

• Core Methodology: Profiling provides the data; bottleneck analysis provides the diagnosis.
• Common Bottlenecks: Memory bandwidth, latency of dependent operations, compute throughput, or host-device data transfer.
• Outcome: Guides optimization efforts. For example, if instrumentation profiling reveals a kernel is memory-bound, optimizations like memory coalescing or improved cache hit rate are pursued.