Papers
Topics
Authors
Recent
Search
2000 character limit reached

Latency-Based Tiling: Memory Optimization

Updated 13 July 2026
  • Latency Based Tiling is a systems approach that leverages measured latency to determine optimal tile shapes and improve memory hierarchy utilization.
  • It uses triangular loop patterns and slope analysis of miss ratio scaling to identify cache boundaries with minimal compile time overhead.
  • The method is applicable across CPUs, GPUs, and FPGAs, balancing tile management overhead with gains in data locality and performance.

Latency Based Tiling denotes a systems based approach to deriving approximate tiling solution that maximizes locality while maintaining a fast compile time. In its canonical formulation, the method uses triangular loops to characterize miss ratio scaling of a machine, treats sharp increases in data-access latency as indicators that the data footprint exceeds a cache level, and then tiles affine loop nestings in a subset of the polyhedral model based on both the derived memory hierarchy and the observed data footprint per iteration (Cashman, 26 Sep 2025). In the broader literature, closely related latency-oriented tiling schemes appear whenever tile shape, tile placement, or tile schedule are chosen primarily to reduce exposed memory, communication, or I/O latency rather than merely to partition work.

1. Definition and measurement-driven formulation

The paper "Latency Based Tiling" defines the method as a systems based approach that derives an approximate tiling solution by directly measuring latency rather than by running an online autotuning search (Cashman, 26 Sep 2025). Its measurement kernel uses triangular loops to characterize miss ratio scaling while avoiding prefetcher distortion. Miss ratio scaling is used to capture the relationship between data access latency and working set size, with noticeable increases in latency indicating approximate locations for L1, L2, and L3 memory sizes. These sizes are expected to be under approximations of a systems true memory sizes, which the paper explicitly relates to the shared nature of cache in a multi process system and to defensive loop tiling. The implementation is in Rust, is described as hardware agnostic, and is intended to be portable and memory safe.

The method computes a discrete slope of the latency curve,

ΔtΔs=ti+1tisi+1si,\frac{\Delta t}{\Delta s} = \frac{t_{i+1}-t_i}{s_{i+1}-s_i},

and uses the largest derivative peaks as cache-boundary candidates (Cashman, 26 Sep 2025). That formulation makes latency itself the observable used to infer tile-relevant cache capacities. In contrast to autotuning, which the paper describes as effective but prohibitively slow, the stated objective is negligible compile time overhead.

Empirically, the reported PolyBench results are mixed but favorable on memory-intensive kernels: 3mm reaches a speedup of 1.21×1.21\times, gemm reaches 1.15×1.15\times, and 2mm reaches 1.08×1.08\times, while covariance and jacobi-2d slightly regress to 0.98×0.98\times and 0.97×0.97\times, respectively (Cashman, 26 Sep 2025). This establishes a central property of latency-based tiling: it is not a universal accelerator for all loop nests, but a locality-driven transformation whose gains depend on whether cache-latency reduction outweighs tile-management overhead.

2. Cache locality, reuse distance, and compiler-level tiling

Several compiler systems formulate the same objective in different terms. Parakeet introduces TiledMap, TiledReduce, and TiledScan, which break computations into local pieces of bounded size so as better to make use of small fast memories (Hielscher et al., 2013). Applying the transformation once tiles for cache, and applying it again enables tiling for registers. Cache-tile sizes are left unspecified until runtime, when an autotuning search is performed; register tiles are fixed at compile time and then fully unrolled so that scalar replacement can place the data in registers. The paper reports significant speedups on locality-sensitive benchmarks and states that register tiling accounted for 60% of the speedups due to tiling on average (Hielscher et al., 2013).

A more explicitly cache-theoretic line of work appears in "Model-Driven Automatic Tiling with Cache Associativity Lattices" (Adjiashvili et al., 2015). That paper argues that conflict misses are the only fundamentally important cache miss category, models cache behavior through associativity lattices, and claims two theoretical advantages over rectangular tiles: volume maximization and miss regularity. The implementation reports speedups of over 10×10\times versus unoptimized code and competitiveness with aggressive GCC and ICC optimizations, even though it currently tiles for only one cache level (Adjiashvili et al., 2015). The resulting perspective is more structural than capacity-based: latency is minimized by constraining per-set conflict structure and reuse distance, not simply by fitting an estimated working set into total cache capacity.

Stencil compilers expose a related dependence-sensitive variant. In Devito, time tiling requires skewing so that transformed dependence vectors are lexicographically positive, after which space-time tiles contain multiple iterations in the outermost time dimension and improve data reuse in all dimensions (McCormick, 2017). On the acoustic wave equation, the reported runtime improvement reaches 27.5%. BoxLib embeds tiling directly into its AMR framework through MFIter, so that application codes can switch from grid-level iteration to tile-level iteration without changing the Fortran kernel interface, and the paper presents this as a path to effective thread utilization on architectures with many more cores per node (Zhang et al., 2016). Taken together, these systems show that latency-based tiling can be driven by runtime search, by cache-associativity models, or by polyhedral legality constraints, but the operative goal remains the same: make reuse happen before cache eviction.

3. GPU latency hiding, tile geometry, and memory-layout co-design

On GPUs, latency-based tiling is typically expressed through occupancy, coalescing, synchronization cost, and overlap of memory transfer with computation. "Tiling for Performance Tuning on Different Models of GPUs" studies block-level tiling for bilinear image interpolation and shows that the same tile is not uniformly optimal across architectures (Xu et al., 2010). The paper compares GTX 260 and GeForce 8800 GTS, emphasizes that global memory latency dominates, and reports that tile 32×432\times4 performs best at higher scales on both GPUs. It also explicitly states that an optimized tiling strategy on one GPU model is not always a good solution on another.

Warp-specialized image-processing work sharpens that viewpoint. "Model-Based Warp Overlapped Tiling for Image Processing Programs on GPUs" fuses loops to form overlapped tiles that fit in a single warp, uses lightweight warp synchronization in place of __syncthreads(), and introduces hybrid tiling that stores overlapped regions in a combination of thread-local registers and shared memory (Jangda et al., 2019). The reported result is 1.65×1.65\times speedup over Halide’s manual schedules on an NVIDIA GTX 1080Ti and 1.33×1.33\times on an NVIDIA Tesla V100. The technical point is that tile size is chosen jointly with synchronization scope and on-chip storage placement, so latency hiding depends on warp-level execution geometry rather than only on data reuse.

Tile-oriented GPU DSLs internalize the same scheduling logic. TileLang presents a generalized tiled programming model in which tiles move between DRAM and SRAM and then undergo a sequence of computations, while scheduling space—thread binding, layout, tensorize, and pipeline—is decoupled from dataflow (Wang et al., 24 Apr 2025). The paper attributes state-of-the-art performance in key kernels to this unified block-and-thread paradigm and to transparent scheduling capabilities. TileLens addresses a different latency source: large-granularity memory systems with kilobyte-scale accesses (Ju et al., 4 Jul 2026). It identifies read amplification caused by a mismatch between two-dimensional compute tiles and one-dimensional memory layout, proposes tile-major layout, and reports that combining tile-major layout with an adaptive hardware prefetcher reduces the geomean slowdown from 1.21×1.21\times0-1.21×1.21\times1 with conventional layouts to within 1% of an HBM-only baseline on HBF-augmented GPUs with a 1.21×1.21\times2 HBF NAND read latency (Ju et al., 4 Jul 2026). TileMaxSim then shows an I/O-aware form of GPU tiling for retrieval scoring: multi-query SRAM tiling, dimension tiling, and fused product-quantization scoring reach 80.2% of peak HBM bandwidth and reduce ColBERT scoring latency at 100K candidates from 268 ms to 1.2 ms (Sharma, 24 Jun 2026). These results collectively make clear that, on GPUs, latency-based tiling is as much about tile-aligned memory layout and synchronization structure as about tile size itself.

4. Memory-constrained edge inference and FPGA-oriented formulations

Latency-based tiling also appears when the limiting resource is memory footprint rather than raw arithmetic throughput. MAFAT extends prior distributed partitioning work into a memory-aware execution strategy for a single device and allows multiple groups of convolutional layers to be fused and tiled independently (Farley et al., 2021). The paper proposes a memory usage predictor and a search algorithm for optimized fusing and tiling configurations, and reports that on YOLOv2 the method can run in less than half the memory, with a speedup of up to 2.78 under severe memory constraints; the returned configuration is within 6% of the best latency measured in a manual search (Farley et al., 2021). Here the tile is not primarily a cache-fitting object. It is a memory-budgeting object chosen to avoid swapping while keeping overhead low.

The FPGA literature frames the same issue through dataflow and hierarchy. A recent survey on edge-AI FPGA accelerators describes tiling or blocking as partitioning large tensors so they fit into on-chip memory tiers, and argues that latency is dominated by how computation is tiled and mapped onto the memory hierarchy and PE array (Li, 13 May 2025). It organizes the discussion around Weight-Stationary, Output-Stationary, Row-Stationary, and No-Local-Reuse dataflows; reviews MAESTRO and Timeloop as analytical frameworks; and emphasizes multi-level tiling plus loop unrolling and pipelining. This suggests that, in FPGA settings, latency-based tiling is inseparable from the chosen dataflow: a tile is useful only insofar as it creates reuse at the memory tier that the architecture can exploit.

5. Spatial dataflow, communication tolerance, optical systems, and MoE inference

Spatial dataflow accelerators reinterpret tiling as a mapping problem over distributed cores and on-chip networks. TL compiles tile-based programs, such as Triton kernels, onto spatial dataflow architectures and focuses on distributing tile instances across spatially distributed cores while exploiting the on-chip network and distributed memories to increase data reuse and reduce communications (Li et al., 17 Dec 2025). Its analytical performance model explicitly combines tile compute time, data movement time, and overlap through a double-buffered pipeline, and the reported evaluation shows geomean speedup of about 1.21×1.21\times3 over Tenstorrent’s TTNN library, with up to 1.21×1.21\times4 speedups on some GEMM shapes (Li et al., 17 Dec 2025). In this setting, latency-based tiling is fundamentally spatiotemporal: tile placement, broadcast pattern, and buffering policy determine whether communication remains on-chip or spills to high-latency global memory.

Task-graph formulations push the same logic to distributed execution. "Task Graph Transformations for Latency Tolerance" derives latency-tolerant execution by partitioning tasks into sets often denoted 1.21×1.21\times5, 1.21×1.21\times6, and 1.21×1.21\times7, so that communication can overlap with local work under strong scaling (Eijkhout, 2018). A plausible implication is that tile boundaries here are not geometric blocks but macro-steps in a task graph whose size is chosen to amortize latency and tolerate communication delay.

Optical accelerators show that the concept is not limited to electronic memory hierarchies. "Channel Tiling for Improved Performance and Accuracy of Optical Neural Network Accelerators" proposes a channel tiling scheme for 4F CNN systems that uses the high resolution of the 4F system to perform channel summation inherently in the optical domain before sensor detection (Li et al., 2020). Compared to the state of the art, the paper reports similar accuracy, 33% improvement in required sensing precision, 10dB reduction in tolerable sensing noise, 1.21×1.21\times8 total filters required, 1.21×1.21\times9-1.15×1.15\times0 throughput improvement, and as much as 1.15×1.15\times1 reduction in required output camera resolution or bandwidth, without any additional optical hardware (Li et al., 2020). TileQ brings the idea back to expert-based LLMs: it uses 2D-tiling structured low-rank quantization for MoE experts and introduces a fused inference technique that reduces inference latency to approximately 5% while cutting additional memory usage up to 1.15×1.15\times2 (Gu et al., 10 May 2026). These systems differ radically in substrate, but each uses tiling to move a dominant reduction or transfer into a faster domain.

6. Trade-offs, misconceptions, and research directions

A recurring misconception is that latency-based tiling is simply “choose the largest tile that fits.” The literature consistently rejects that simplification. On GPUs, a tile that maximizes occupancy on GTX 260 can underutilize GeForce 8800 GTS because per-SM limits change the relationship between block size and latency hiding (Xu et al., 2010). On stencils, time tiling is not legal until skewing makes dependence vectors lexicographically positive, and buffered time dimensions impose additional tile-size constraints (McCormick, 2017). On sparse CNN activations, very fine-grained compact tiles approach ideal data selectivity but can incur 25% indexing overhead, whereas GrateTile’s structured uneven subtensors reach about 55% DRAM bandwidth reduction with only 0.6% of feature map size for indexing storage (Lin et al., 2020). On large-granularity memory systems, even high-bandwidth devices can be slowed down by up to an order of magnitude if tile layout is mismatched to access granularity, because read amplification dominates (Ju et al., 4 Jul 2026).

Another misconception is that latency-based tiling obviates autotuning. The record is more nuanced. Parakeet leaves cache-tile sizes unspecified until runtime and searches for them online (Hielscher et al., 2013), whereas "Latency Based Tiling" explicitly seeks negligible compile time overhead by replacing search with direct latency characterization (Cashman, 26 Sep 2025). This suggests a spectrum rather than a dichotomy: empirical search, analytical modeling, direct latency probing, and domain-specific scheduling can all serve the same objective, but with different compile-time and portability costs.

Current research directions are correspondingly diverse. The FPGA survey highlights hybrid dataflows, partial reconfiguration, and compiler flows as future directions for edge AI (Li, 13 May 2025). TL emphasizes richer performance models and broader back-end support for spatial dataflow fabrics (Li et al., 17 Dec 2025). TileLang presents a programming model in which pipeline, layout, and thread binding remain explicit enough for compilers to synthesize latency-hiding schedules (Wang et al., 24 Apr 2025). The cumulative picture is that latency-based tiling has evolved from cache blocking into a general systems methodology: characterize the dominant latency source, choose a tile shape and traversal that localize it, and place the reduction, communication, or fetch into the fastest domain the hardware exposes.

Topic to Video (Beta)

No one has generated a video about this topic yet.

Whiteboard

No one has generated a whiteboard explanation for this topic yet.

Follow Topic

Get notified by email when new papers are published related to Latency Based Tiling.