Configuration Space

The dimensionality of a configuration space is defined by the number of independent, tunable parameters. Each parameter adds a new axis to the search space. For example, a kernel's configuration space might be defined by:

• Workgroup Size (X, Y dimensions)
• Tile Size for local memory
• Loop Unroll Factor
• Vectorization Width High dimensionality leads to an exponential increase in possible combinations, making exhaustive search infeasible and necessitating intelligent search algorithms like Bayesian Optimization.

Each tunable parameter has a specific type and a defined domain of valid values. These domains constrain the search.

• Integer Parameters: e.g., threads_per_block ∈ {32, 64, 128, 256, 512}
• Categorical Parameters: e.g., memory_scheme ∈ {shared, global, texture}
• Ordinal Parameters: e.g., loop_order where (i,j,k) is distinct from (k,j,i) The domain can be continuous (e.g., a learning rate) or discrete. For NPU kernel tuning, parameters are almost always discrete and constrained by hardware limits (e.g., maximum threads per block).

Configuration spaces for hardware optimization are fundamentally discrete and combinatorial. The total number of configurations is the product of the domain sizes for each parameter. For a kernel with 5 parameters, each with 10 possible values, the space contains 100,000 distinct points. This combinatorial explosion is the primary challenge for auto-tuning, as evaluating every point via direct profiling is computationally prohibitive, requiring heuristic or model-guided search strategies.

The configuration space maps to a performance landscape, where each point (configuration) has an associated cost metric (e.g., execution time, power consumption). This landscape is:

• Non-linear: Small parameter changes can cause large, discontinuous performance shifts.
• Non-convex: Contains many local minima and maxima.
• Noisy: Repeated measurements of the same point can vary due to system noise. Understanding this rugged terrain is why simple search methods like grid search are inefficient compared to techniques that build a surrogate performance model.

Not all parameter combinations are valid or legal. Constraints define invalid regions of the space that must be pruned during search. Common constraints in NPU tuning include:

• Resource Limits: Total shared memory usage per block cannot exceed hardware capacity.
• Divisibility Requirements: Tile sizes must evenly divide problem dimensions.
• Hardware Alignment: Memory access patterns must adhere to byte-alignment rules for coalescing. An effective auto-tuner must either sample only from the valid subspace or have a mechanism to penalize invalid configurations heavily.

Parameters exhibit varying degrees of sensitivity (how much they affect performance) and correlation (how they interact).

• A highly sensitive parameter, like workgroup_size, causes large performance swings.
• Correlated parameters interact; the optimal value for one depends on the value of another (e.g., optimal tile_size depends on chosen memory_scheme). Identifying these relationships allows search algorithms to focus on important dimensions and model interactions, drastically reducing the number of profiled samples needed to find a near-optimal configuration.

The automated process of searching a configuration space to find the parameter set (e.g., tile size, workgroup dimensions) that yields optimal performance for a computational kernel on specific hardware. It replaces manual trial-and-error with systematic search algorithms.

• Core Mechanism: Iteratively evaluates kernel variants defined by different parameter combinations.
• Objective: Maximize metrics like compute throughput or minimize latency.
• Methods: Can use exhaustive search, random sampling, or guided methods like Bayesian Optimization.

The systematic exploration of a configuration space. This is the algorithmic heart of auto-tuning, defining how the space is navigated.

• Search Strategies: Range from brute-force grid searches to sophisticated model-based approaches.
• Challenge: The space is often vast and non-convex; efficient search is critical.
• Goal: To find a global or high-quality local optimum for performance without evaluating every possible point.
• Example: Searching over workgroup_size = {32, 64, 128, 256} and vectorization_factor = {2, 4, 8}.

An analytical or machine-learned model that predicts execution time or resource usage of a kernel based on its parameters and hardware specs. It acts as a surrogate for expensive physical measurements during parameter search.

• Purpose: Guides auto-tuning by predicting the outcome of untested configurations.
• Types: Can be simple analytical models (e.g., roofline model) or complex learned models (e.g., neural networks).
• Inputs: Configuration parameters, hardware characteristics (peak FLOPS, memory bandwidth), and problem size.
• Output: Predicted runtime, helping prioritize promising regions of the configuration space.

A software tool that measures the actual runtime behavior of a kernel executing on hardware. It provides the ground-truth data that defines the performance landscape of the configuration space.

• Measures: Execution time, compute throughput, memory bandwidth utilization, cache hit rate, and hardware performance counter events.
• Role in Tuning: Used to evaluate specific points in the configuration space, feeding data back to the auto-tuner or performance model.
• Types: Includes sampling profilers (low overhead) and instrumentation profilers (high detail).

The process of identifying the primary limiting factor that restricts a kernel's performance. It determines why a point in the configuration space performs poorly, guiding the search for better configurations.

• Key Classifications:
  • Compute Bound: Limited by ALU speed; improve by increasing occupancy or vectorization factor.
  • Memory Bound: Limited by memory bandwidth; improve via memory coalescing or better cache use.
  • Latency Bound: Limited by dependency stalls or pipeline stalls.
• Tool Use: Relies on data from kernel profilers and performance counters.

A sequential, model-based parameter search strategy. It is particularly effective for optimizing expensive-to-evaluate functions, like kernel runtimes, where the configuration space is large.

• Mechanism: Builds a probabilistic surrogate model (often a Gaussian Process) of the performance landscape and uses an acquisition function to decide the most informative configuration to test next.
• Advantage: Intelligently balances exploration (testing uncertain regions) and exploitation (refining known good regions).
• Outcome: Typically finds a near-optimal configuration with far fewer evaluations than random or grid search.