Tiled High-order Sliding Window Attention

• Tiled high-order sliding window attention is a local mechanism that groups tokens into tiles and slides windows over N-dimensional data.
• It enhances efficiency by aligning computations with hardware, reducing irregular memory patterns and accelerating long-sequence processing.
• It plays a key role in hybrid architectures by complementing global linear recurrence, achieving competitive accuracy with reduced latency in image and video tasks.

Tiled high-order sliding window attention is a local attention mechanism for high-order, or $N$-dimensional, data in which attention is restricted to local windows across multiple dimensions and implemented through tile-aligned computation rather than token-by-token masking. In the ENA line of work, the mechanism is identified as tiled high-order sliding window attention and implemented as Sliding Tile Attention (STA): tokens are grouped into tiles, all tokens in a tile share the same local attention window, and the window slides tile-by-tile across the input. The resulting design is presented as efficient in both theory and practice for long sequences of high-order data, especially when combined with linear recurrence in hybrid architectures (Zhong, 16 Aug 2025).

1. Definition and scope

Sliding window attention (SWA) is a local attention mechanism in which each token attends only to tokens within a fixed-size local neighborhood rather than to the full sequence. When generalized to high-order data, such as images and videos, SWA operates over local windows in multiple dimensions rather than only along a 1D temporal axis. Tiled high-order SWA refines this construction by grouping adjacent tokens into tiles and shifting the attention window tile-by-tile, so that all tokens in a tile share the same local window (Zhong, 16 Aug 2025).

In ENA, this tiled form is described more precisely as Sliding Tile Attention. Its defining properties are that it is a natural extension of 1D sliding window attention to high-dimensional data, that it imposes no boundary restrictions for queries unlike naive block attention, and that its tiled design avoids inefficient memory patterns by eliminating mixed or partial blocks. The same basic distinction is emphasized in video generation work on STA, where token-wise SWA is contrasted with tile-wise attention in 2D and 3D settings: all queries within a tile attend to the same set of key tiles, so the attention map contains only dense blocks and empty blocks, while mixed blocks are eliminated (Zhang et al., 6 Feb 2025).

Within this literature, “high-order” refers to $N$-dimensional data layouts such as $N=2$ for images and $N=3$ for video. The tiled variant is therefore not merely a sparse mask on a flattened sequence. It is a locality mechanism defined over the original spatial or spatiotemporal structure and then organized to align with hardware execution.

2. Mathematical formulation

For $N$-dimensional input, ENA writes the input tensor as

$$\mathbf{X} \in \mathbb{R}^{B \times L_1 \times L_2 \times \cdots \times L_N \times D},$$

with attention window size $W_1 \times W_2 \times \cdots \times W_N$ and tile size $T_1 \times T_2 \times \cdots \times T_N$. If ${\bm i} = (i_1, i_2, \ldots, i_N)$ denotes a tile index, then every tile ${\bm i}$ is associated with a shared attention window $N$0. For a position $N$1 inside tile $N$2, the output is

$N$3

where $N$4 are attention weights computed within the local window. In practice, the window slides tile-by-tile, so the stride equals the tile size (Zhong, 16 Aug 2025).

For 3D video attention, STA work states the same idea in a block-sparse form. The $N$5 token grid is partitioned into non-overlapping tiles of size $N$6, and for each query tile all queries attend to the same set of key tiles within a surrounding window of size $N$7. If both $N$8 and $N$9 are multiples of $N=2$0, the number of dense blocks per query tile is $N=2$1, and the total number of dense blocks is

$N=2$2

This is the formal expression of the tile-sharing rule: sparsity is defined at block granularity rather than by per-token masking (Zhang et al., 6 Feb 2025).

A central mathematical consequence is that tiled high-order SWA preserves local attention semantics while changing the execution unit. Queries remain local in the original $N=2$3-dimensional domain, but computation is scheduled over aligned tiles and dense submatrices.

3. Tiling, sliding, and hardware alignment

The implementation sequence in ENA is explicit. First, the input tensor is partitioned into non-overlapping tiles across all spatial or spatiotemporal dimensions. Second, the window moves with stride equal to the tile size. Third, all tokens in a tile attend to a fixed set of neighboring tokens within the associated local window. Fourth, the attention computation is organized to maximize memory reuse and minimize overhead. Fifth, unlike block attention, there are no boundary restrictions for queries: any token can attend up to the window boundary regardless of tile or spatial boundary (Zhong, 16 Aug 2025).

The hardware rationale is that coarse-grained sliding ensures that no attention block is mixed. ENA states that this avoids inefficient memory patterns and enables real speedups rather than merely theoretical FLOPs reduction. The video-generation STA paper sharpens the same point: SWA in 2D and 3D leads to highly irregular and mixed attention blocks on the GPU, whereas STA restructures attention so that the matrix contains only dense blocks, which are fully computed, and empty blocks, which are fully skipped (Zhang et al., 6 Feb 2025).

Kernel-level reports reinforce that this is an implementation principle rather than a cosmetic reformulation. The STA implementation for video generation is built on top of FlashAttention3 and ThunderKittens, uses a producer-consumer paradigm, asynchronously loads only necessary key and value tiles, and skips empty blocks entirely. Under these optimizations, STA achieves 58.79% MFU, accelerates attention by 2.8-17x over FlashAttention-2 and 1.6-10x over FlashAttention-3, and on HunyuanVideo reduces end-to-end latency from 945s to 685s without quality degradation; with finetuning, latency is reduced to 268s with only a 0.09% drop on VBench (Zhang et al., 6 Feb 2025).

This hardware-aligned view also explains why tiled high-order SWA is distinguished from naive sparse attention. The crucial object is not just a local mask, but a sparse pattern whose nonzeros can be grouped into dense, regularly scheduled blocks.

4. Role in hybrid architectures

ENA places tiled high-order SWA inside an alternating hybrid architecture of linear recurrence and local attention. The block-wise update is written as

$N=2$4

$N=2$5

where $N=2$6 alternates between a linear recurrent module and a local attention module. Odd blocks use linear recurrence for global aggregation, and even blocks use local attention for fine-grained dependencies. For short sequences, local attention is replaced with full attention; for long sequences, high-order SWA or STA is used (Zhong, 16 Aug 2025).

The stated intuition is that linear recurrence compresses global information into a state, while SWA complements it by enforcing strict local modeling. This division of labor is central to ENA’s characterization of tiled high-order SWA: it is not introduced as a full replacement for all long-range reasoning, but as the local branch in a hybrid system for ultra-long high-order data (Zhong, 16 Aug 2025).

Subsequent work on hybrid allocation further situates tiled SWA as one point in a larger design space. ConSA learns FA/SWA assignment under a user-specified sparsity target and explicitly notes that the framework is directly applicable to any pair or set of sparse attention mechanisms, including tiled SWA and higher-dimensional windows. In that setting, FA/SWA can be assigned per layer or per KV head, with learned allocations concentrating FA into contiguous middle-layer blocks and placing SWA in the bottom layers (Chen et al., 16 Jun 2026).

A plausible implication is that tiled high-order SWA is most naturally understood as a modular local operator whose value depends on how it is paired with global mechanisms, rather than as an isolated alternative to full attention.

5. Empirical behavior and comparative performance

ENA reports that tiled high-order SWA is effective across image classification, video understanding, and image and video generation tasks. For long sequences, ENA with 2D STA using tile size $N=2$7 and a window covering approximately 30% of tokens, or 70% sparsity, outpaces Transformer+FlashAttention in both speed and memory. The same study states that accuracy saturates with approximately 30% window coverage and that wider windows yield diminishing gains (Zhong, 16 Aug 2025).

Selected quantitative examples are summarized below.

Setting	Full / baseline	Tiled high-order SWA result
ImageNet-1K, seq=4096	DeltaNet + Full Attention: 66.11% (14h42m)	DeltaNet + 2D STA: 66.41% (10h55m)
ImageNet 512×512 generation	Transformer-based baselines	ENA with STA: FID = 4.87
HunyuanVideo, 5s 720p	FA3: 945s	STA: 685s; finetuned STA: 268s

These measurements support two recurring claims in the literature. First, tiled high-order SWA can match or exceed full attention in accuracy while improving runtime in long-sequence regimes. Second, the benefit is especially visible when the attention operator is implemented as a tile-aligned sparse kernel rather than as a token-wise masked dense kernel (Zhong, 16 Aug 2025, Zhang et al., 6 Feb 2025).

The comparative evidence against block attention is also consistent. ENA states that SWA or STA outperforms block attention mainly because block attention suffers from boundary effects and limits expressivity near block borders, whereas SWA or STA does not have this problem. The video-generation STA work makes the same distinction in hardware terms: Swin is efficient because it sacrifices expressiveness by non-overlapping, non-sliding windows, and this degrades quality (Zhong, 16 Aug 2025, Zhang et al., 6 Feb 2025).

6. Misconceptions, related variants, and open directions

One common misconception is that any local attention over blocks is equivalent to tiled high-order SWA. The cited work rejects this equivalence. Block attention imposes non-overlapping local regions and introduces boundary effects; tiled high-order SWA preserves sliding locality and removes query-side boundary restrictions. Another misconception is that sparse attention with favorable asymptotic FLOPs automatically yields practical speedups. ENA explicitly states that kernel-level optimization is still evolving and that its current Triton-based implementation is less optimized than pure CUDA implementations such as FlashAttention v2, with further speedup possible through custom kernels such as ThunderKittens and CUDA (Zhong, 16 Aug 2025).

The broader SWA literature also shows that local attention is not a single mechanism. Reference Sliding Window Attention in OCR fixes a reference segment and lets only the output window slide, maintaining constant KV cache throughout decoding. Stochastic Attention applies a random permutation before windowed attention and restores the original order afterward, turning a fixed local window into a stochastic global one within the same $N=2$8 per-layer budget. These are not tiled high-order SWA in the ENA sense, but they demonstrate that the sliding-window paradigm is being reworked along orthogonal axes: persistent reference access, stochastic routing, and hardware alignment (Yin et al., 22 Jun 2026, Jin et al., 1 Apr 2026).

Architecture-level studies add a final caution. Work on hybrid LLMs reports that SWA window size primarily affects how fast long-context capability emerges and identifies a phenomenon called Large-Window Laziness, in which larger SWA windows can delay the formation of retrieval heads in full-attention layers. Work on math reasoning similarly shows that SWA may lag behind full self-attention after supervised conversion, with reinforcement learning narrowing the gap by adapting trajectories to the SWA constraint (Qiao et al., 13 Jun 2026, Liu et al., 10 Jun 2026).

Taken together, these results define the current status of tiled high-order SWA with some precision. It is a hardware-aligned, high-dimensional local attention mechanism whose principal advantages arise when locality, tile granularity, and kernel design are co-designed. Its strongest empirical role is in hybrid systems that pair strict local modeling with a complementary global mechanism, and its main open problems remain kernel optimization, broader modality coverage, and architecture-specific allocation of local versus global attention (Zhong, 16 Aug 2025).