Memory Bound

The defining characteristic of a memory-bound kernel is its low arithmetic intensity, measured in FLOPS/byte. This ratio quantifies the number of floating-point operations performed per byte of data transferred from main memory.

• Example: A simple vector addition (A[i] = B[i] + C[i]) performs only 1-2 operations per 12-24 bytes loaded/stored, resulting in very low arithmetic intensity.
• Impact: The theoretical peak compute throughput (e.g., TFLOPS) of the NPU cannot be achieved because the memory subsystem cannot supply data fast enough to keep the compute units busy.

A clear indicator of a memory-bound state is when the measured memory bandwidth approaches the hardware's theoretical peak bandwidth. Profiling tools will show:

• Memory controller utilization at or near 100%.
• Compute unit utilization (ALU occupancy) significantly lower, often below 50-70%.
• Performance counters revealing a high number of memory transactions but a low count of issued compute instructions.

When bandwidth is saturated, issuing more concurrent memory requests does not improve performance and can increase latency due to contention.

Memory-bound workloads often exhibit poor cache locality, leading to a low cache hit rate. This forces excessive accesses to slower, higher-latency main memory (DRAM).

• Symptoms: High rates of L1/L2 cache misses and DRAM accesses per instruction.
• Cause: Data access patterns are often strided, random, or use large working sets that exceed cache capacity.
• Consequence: The memory access latency, not the compute operation latency, becomes the dominant factor in the kernel's execution time. Techniques like tiling and memory coalescing aim to improve this.

Specific computational patterns are inherently prone to being memory-bound on NPUs and GPUs:

• Element-wise operations: Activation functions (ReLU, Sigmoid), tensor addition.
• Reductions: Summing or finding the max/min across a large array.
• Bandwidth-limited layers: The input/output operations of fully-connected (dense) layers and some embedding lookups.
• Data shuffling: Transposes, gathers, scatters, and certain data format conversions.
• Stencil operations with large footprints that exceed cache capacity.

These kernels perform minimal computation on each data element loaded, making performance directly proportional to memory bandwidth.

Understanding the opposite condition clarifies the memory-bound state.

Identifying a memory-bound bottleneck requires specific profiling metrics:

• Essential Metrics:
  • dram__bytes_read.sum & dram__bytes_write.sum: Total DRAM traffic.
  • sm__throughput.avg.pct_of_peak_sustained_elapsed: Compute utilization.
  • l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum: Global load transactions.
• Key Ratios:
  • Achieved Bandwidth / Peak Bandwidth: Ratio > 0.7-0.8 suggests bandwidth saturation.
  • Arithmetic Intensity: Calculate from FLOPS and bytes accessed.
• Tool Use: NPU vendor profilers (e.g., NVIDIA Nsight Compute, AMD ROCprof) and performance models are used to compare measured metrics against hardware limits to confirm the bottleneck.

Memory coalescing is a critical optimization where consecutive threads in a warp or wavefront are programmed to access consecutive memory addresses. This pattern allows the memory controller to combine multiple scattered requests into a single, wider transaction, dramatically increasing effective bandwidth.

• Mechanism: Instead of 32 separate 4-byte loads, a coalesced access can become one 128-byte cache line fetch.
• Impact: Can improve memory throughput by an order of magnitude for structured data like matrices and tensors.
• Implementation: Requires careful data layout (e.g., Structure of Arrays vs. Array of Structures) and thread indexing.

Loop tiling partitions large data arrays into smaller blocks (tiles) that fit into the NPU's fast, on-chip memory (shared memory/L1 cache). This transforms a memory-bound problem by reusing loaded data multiple times for computation.

• Goal: Maximize data locality to reduce accesses to slower global memory.
• Key Parameter: Tile size selection is auto-tuned based on cache size, register count, and kernel resource usage.
• Example: In a matrix multiplication, tiles of matrices A and B are loaded once, then used for multiple partial sum calculations.

Prefetching is a technique that proactively loads data into a cache or registers before it is explicitly needed by the compute units, hiding memory access latency. It turns serialized load-compute stalls into overlapped operations.

• Software Prefetching: Explicit compiler intrinsics or directives (e.g., __prefetch) hint to the memory subsystem.
• Hardware Prefetching: NPU hardware automatically detects sequential access patterns and fetches ahead.
• Use Case: Essential for kernels with predictable, strided access patterns to keep the computational pipeline fed.

NPUs feature a small, low-latency shared memory (scratchpad) that is explicitly managed by the programmer. For memory-bound kernels, this memory acts as a manually controlled cache for frequently reused data.

• Process: Data is loaded from global memory into shared memory cooperatively by a thread block, then accessed at high bandwidth by all threads in the block.
• Optimization: Eliminates redundant global memory loads and enables efficient inter-thread communication.
• Challenge: Requires careful synchronization (__syncthreads()) and awareness of bank conflicts to avoid serialization.

Kernel fusion combines multiple consecutive computational kernels into a single kernel launch. This optimization is powerful for memory-bound workloads because it eliminates intermediate results being written to and read from global memory.

• Benefit: Transforms memory-bound intermediate steps into compute-bound fused operations, staying within registers or cache.
• Example: Fusing an activation function (e.g., ReLU) with a preceding convolution layer, so the convolution output is activated before being stored.
• Trade-off: Increases register pressure and kernel complexity, which must be balanced.

Utilizing the NPU's vector load/store instructions (e.g., loading 128-bit vectors) maximizes the utilization of the memory bus per transaction. This reduces the instruction overhead and number of memory requests issued by the threads.

• Mechanism: A single thread uses a wide data type (e.g., float4) or intrinsic to access multiple data elements in one instruction.
• Requirement: Data must be aligned in memory (often to 16-byte boundaries) for peak performance.
• Auto-Tuning: The vectorization factor (2, 4, 8) is a key parameter searched during auto-tuning to match the hardware's optimal memory transaction size.

The opposite state of a memory-bound workload. A program is compute-bound when its execution time is limited by the speed of the arithmetic logic units (ALUs), not by memory or I/O speeds. In this state, the processor's computational pipelines are fully saturated, and increasing memory bandwidth yields no performance improvement.

• Key Indicator: High compute throughput (e.g., FLOPS) nearing the hardware's peak, but low memory bandwidth utilization.
• Optimization Focus: Shifts to improving instruction-level parallelism, reducing arithmetic operation latency, and maximizing ALU occupancy.

The maximum theoretical rate at which data can be read from or written to a memory subsystem by a processor, measured in gigabytes per second (GB/s). This is the primary resource that a memory-bound kernel is exhausting.

• Critical Metric: The actual achieved bandwidth versus the hardware's peak bandwidth indicates the severity of the memory bottleneck.
• Factors: Determined by memory bus width, clock speed, and the efficiency of access patterns (e.g., memory coalescing).
• Tooling: Measured directly using hardware performance counters or inferred from kernel profiler data.

A critical memory access pattern optimization for parallel architectures like NPUs and GPUs. It occurs when consecutive threads in a warp or wavefront access consecutive memory locations, allowing the memory controller to combine these accesses into a single, wide transaction.

• Impact on Memory-Bound Kernels: Poor coalescing (stride or random access) drastically reduces effective memory bandwidth, exacerbating the memory-bound condition.
• Optimization Technique: Restructuring data layouts (e.g., Array of Structures to Structure of Arrays) and thread indexing to ensure contiguous, aligned access.

The percentage of memory access requests that are successfully served from a fast cache memory (L1, L2), avoiding a slower access to main memory (DRAM). A low cache hit rate is a strong contributor to a memory-bound state.

• Performance Relationship: High hit rates reduce the demand on main memory bandwidth, alleviating memory-bound pressures.
• Optimization Strategies: Increasing temporal locality (reusing data) and spatial locality (accessing nearby data) through techniques like loop tiling (tile size selection) and blocking.

The systematic process of identifying the primary limiting factor that restricts the overall performance of a system or application. Diagnosing a memory-bound kernel is a key outcome of this analysis.

• Methodology: Uses data from kernel profilers, execution traces, and performance counters to compare resource utilization (compute vs. memory).
• Roofline Model: A visual analytical model that plots algorithmic performance (FLOPS) against operational intensity, clearly showing if a kernel is memory-bound or compute-bound.
• Goal: To direct optimization efforts to the correct subsystem, avoiding wasted effort on non-critical paths.

The automated process of searching a configuration space of kernel parameters to find the optimal setup for a specific hardware target. It is essential for escaping memory-bound conditions.

• Relevant Parameters: Workgroup size, tile size selection, loop unrolling factor, and vectorization factor all dramatically affect memory access patterns and bandwidth utilization.
• Search Techniques: Uses methods like Bayesian optimization or genetic algorithms to efficiently navigate the parameter space.
• Outcome: A kernel tuner automatically discovers configurations that maximize cache usage, improve coalescing, and balance compute/memory workload, directly addressing memory-bound limitations.