Performance Counter

Performance counters are hardware-implemented registers within the processor's execution units, caches, and memory controllers. They count physical events at the silicon level, such as:

• Floating-point operations (FLOPs) executed
• Cache misses (L1, L2, shared memory)
• DRAM read/write transactions
• Branch mispredictions
• Instruction pipeline stalls This provides a ground-truth view of kernel behavior, free from the approximations of software simulation.

Counters track discrete, low-level microarchitectural events rather than high-level software metrics. This allows engineers to correlate performance issues with specific hardware behaviors. For example, a high count of L1 cache misses directly explains memory latency stalls, while a low compute throughput counter indicates the kernel may be memory-bound. This event-driven data is the foundation for constructing accurate performance models used in auto-tuning.

NPUs have a finite number of physical counter registers (often 4-8 per core). Profiling complex kernels may require measuring more events than available hardware counters, necessitating multiplexing—where the profiling tool time-slices events, reducing accuracy. Furthermore, reading counters incurs a small but non-zero overhead. Effective profiling involves strategic selection of the most relevant counters to minimize multiplexing and overhead, focusing on the suspected bottleneck (e.g., memory vs. compute).

Counters can often be configured to measure events for a single core, a subset of cores, or the entire NPU. For parallel workloads, results are typically aggregated across all active threads in a warp or wavefront. Understanding this aggregation is key: a high global cache miss count could be caused by poor access patterns from all threads or severe thread divergence in a few threads. Advanced tools can sometimes provide per-core or per-thread-group insights for finer-grained analysis.

The definitive use of performance counters is to diagnose the root cause of sub-optimal performance. By analyzing ratios and relationships between counters, engineers can classify a kernel's limiting factor:

• Compute-Bound: High FLOPs count, low cache miss rate, low memory bandwidth utilization.
• Memory-Bound: High cache miss rate, memory bandwidth near peak, low FLOP utilization.
• Latency-Bound: High pipeline stall counts due to dependencies or branch mispredictions. This analysis directly informs optimization strategy, such as improving memory coalescing or adjusting the vectorization factor.

These counters measure the utilization of the NPU's arithmetic and logic units (ALUs), which are the core engines for matrix multiplications and convolutions.

• FLOPS (Floating-Point Operations Per Second): The raw rate of FP16, BF16, or FP32 operations executed. A low sustained FLOPS rate often indicates a memory-bound kernel.
• MACs (Multiply-Accumulate Operations): Counts the fundamental operations in dense and convolutional layers. Comparing executed MACs to theoretical peak reveals compute efficiency.
• ALU Utilization Percentage: The fraction of clock cycles where the ALU pipelines are actively processing data, as opposed to stalled. High utilization is the goal for compute-bound workloads.
• Tensor Core/MMA Unit Activity: Specific to NPUs with dedicated matrix-multiply-accumulate units, these counters show usage of the most powerful compute hardware.

Counters in this category track data movement across the NPU's complex memory subsystems, which is often the primary bottleneck.

• Global Memory Transactions: Counts reads/writes to the high-capacity, off-chip DRAM (e.g., HBM). High counts with low compute throughput signal inefficient access patterns.
• Cache Hit/Miss Rates: Measures effectiveness of on-chip SRAM caches (L1, L2, shared memory). A low L1 Cache Hit Rate forces data fetches from slower memory levels.
• Memory Bandwidth Utilization: The percentage of the theoretical peak bandwidth being used. Saturation near 100% confirms a memory-bound condition.
• Bank Conflicts: Occur when multiple threads in a warp attempt to access different data within the same memory bank simultaneously, causing serialized access and reducing effective bandwidth.

These events quantify how effectively the NPU's parallel execution resources are occupied and managed.

• Occupancy: The ratio of active warps/wavefronts to the maximum supported per streaming multiprocessor (SM). Low occupancy can limit ability to hide memory latency.
• Thread Divergence Events: Counts occurrences where threads within a single warp take different execution paths (e.g., in an if/else statement), causing serialization.
• Pipeline Stalls: Tracks cycles where the instruction pipeline is idle due to data dependencies, resource conflicts, or memory wait states.
• Scheduler Issues: Counts cycles where the warp scheduler has no eligible warps to issue, often due to all warps waiting for long-latency operations.

Counters that measure the cost of program control decisions and coordination between threads, which can introduce overhead.

• Branch Instructions & Mispredictions: Counts executed branches and the number that were mispredicted by the hardware, causing pipeline flushes and wasted cycles.
• Barrier Synchronizations: Tracks executions of __syncthreads() or similar instructions. Excessive barriers can serialize execution and reduce parallelism.
• Atomic Operation Contention: Measures conflicts when multiple threads attempt to atomically update the same memory location, causing serialization.
• Predicated Instruction Execution: Counts instructions executed under a false predicate (masked off), representing wasted compute cycles in divergent code.

Events related to transferring data between the host (CPU) and the NPU device, a critical phase for overall pipeline latency.

• PCIe (or Interconnect) Transfer Volumes: Measures the amount of data moved over the host-device link. Overlapping this transfer with computation via asynchronous execution is key.
• DMA Engine Activity: Tracks usage of the Direct Memory Access controllers that manage bulk data transfers without CPU intervention.
• Unified Memory Page Faults: In systems with unified virtual memory, counts page migrations between CPU and NPU memory, which are high-latency events.
• Texture/Surface Memory Accesses: For NPUs supporting these specialized memory spaces, counters track cache efficiency for spatially local data.

Counters that correlate performance activity with energy consumption, crucial for edge and mobile NPU deployments.

• Dynamic Power Events: Estimates power draw from transistor switching activity, often correlated with ALU utilization and clock frequency.
• Memory Access Energy: Models energy consumed by DRAM and cache accesses. Reducing unnecessary global memory transactions directly lowers this.
• Voltage/Frequency Domain Activity: Tracks time spent in different power-performance states (P-states).
• Operations per Joule: A derived metric (e.g., GOPs/J) calculated from compute throughput and power counters, representing the ultimate efficiency metric for the kernel.

A software tool that aggregates and visualizes raw performance counter data, providing a holistic view of kernel execution. It measures execution time, resource utilization, and correlates hardware events to specific lines of source code.

• Integrates multiple counters to identify correlations (e.g., high cache misses with low memory bandwidth).
• Provides timeline views for asynchronous execution and concurrent kernel analysis.
• Essential for hotspot identification, guiding optimization efforts to the most costly code sections.

A chronological, instruction-level record of a program's runtime behavior. While a performance counter aggregates events, a trace provides the exact sequence, crucial for diagnosing complex, non-deterministic issues.

• Captures control flow, memory addresses, and function calls.
• Used to analyze thread divergence and pipeline stalls by seeing the exact path each thread took.
• Heavyweight but precise, often used after counters identify a problematic region for deep inspection.

The systematic process of using profiling data (like counter values) to identify the primary factor limiting performance. It classifies kernels into fundamental categories based on their limiting resource.

• Compute-Bound: Counters show high ALU utilization but low cache miss rates. Limited by arithmetic throughput.
• Memory-Bound: Counters show high DRAM bandwidth usage or high cache miss rates. Limited by data movement speed.
• Latency-Bound: Counters may show low occupancy due to serialized operations or dependency stalls.

An analytical or machine-learned model that predicts execution time or resource usage based on kernel parameters and hardware specs. It uses counter data for validation and calibration.

• Roofline Model is a common analytical model relating operational intensity (Ops/Byte) to attainable performance (FLOPS).
• Guides auto-tuning by predicting the outcome of parameter changes without exhaustive profiling.
• Counters provide ground truth to train and validate data-driven, learned performance models.

The automated process of searching a configuration space to find optimal kernel parameters. It relies on performance counters as the objective function to evaluate each candidate configuration.

• Parameter Search explores variables like workgroup size, tile dimensions, and unroll factors.
• Techniques like Bayesian Optimization use counter feedback (e.g., cache hits, compute throughput) to intelligently navigate the search space.
• Goal is to maximize positive counters (e.g., FLOPS, bandwidth) and minimize negative ones (e.g., stalls, misses).

A performance-degrading condition where multiple executing tasks compete for a shared hardware resource. Performance counters are critical for detecting and quantifying this interference.

• Counters measure shared resource usage like L2 cache access rates, memory controller utilization, and interconnect traffic.
• Manifests in concurrent kernel execution or multi-tenant NPU environments.
• Diagnosis requires comparing counters from isolated vs. concurrent runs to see elevated miss rates or reduced bandwidth.