Optimizing Jagged Flash Attention with TLX for Blackwell GPUs

Deep learning architectures handling dynamic sequence lengths often face a critical bottleneck: memory fragmentation and wasted compute caused by padding. In real-world enterprise workloads—ranging from Meta's Generative Ads Model (GEM) to heterogeneous dynamic batching in modern Large Language Models (LLMs)—input sequences vary drastically in size. Traditional dense tensor representations force developers to pad batches to the longest sequence, leading to severe computational overhead.

To solve this, Jagged Tensors (or nested tensors) eliminate padding by packing contiguous sequences end-to-end. However, implementing hardware-efficient kernels for jagged layouts on next-generation accelerators like the NVIDIA Blackwell B200 GPU introduces complex memory layout alignment and thread synchronization challenges.

In this article, we explore how Jagged Flash Attention (JFA) optimized with Tile Language Extensions (TLX) achieves SOTA execution efficiency on Blackwell, laying the foundation for FlashAttention-4 (FA4). High-performance API aggregation platforms like n1n.ai rely on these underlying low-level kernel optimizations to deliver ultra-low latency inference across modern frontier LLMs.

The Problem with Padded Tensors in Large-Scale AI

When deploying AI models at scale, dynamic batching constructs mini-batches from queries of varying lengths. Consider a batch of four prompt sequences with token lengths of [128, 4096, 512, 2048].

In standard attention implementations, tensors are padded to match the maximum sequence length (N_max = 4096). The resulting matrix requires allocating memory for 4 * 4096 = 16,384 token slots, even though the total valid token count is only 128 + 4096 + 512 + 2048 = 6,784.

This standard padding approach creates two major issues:

1. Wasted Compute: Attention compute complexity scales quadratically $O(N^2)$. Computing Attention matrix multiplications on zero-padded regions wastes up to 60% of GPU TFLOPS.
2. Memory Bandwidth Bottlenecks: Memory bandwidth utilization drops because kernel execution must load zero-padded DRAM regions into High Bandwidth Memory (HBM3e).

Jagged Tensors fix this by storing tokens sequentially in a single flattened 1D/2D continuous buffer, accompanied by an array of cumulative sequence offsets (often denoted as cu_seqlens).

Jagged Tensor Layout:
Data Buffer:  [-- Seq 0 (128) --|-- Seq 1 (4096) --|-- Seq 2 (512) --|-- Seq 3 (2048) --]
Offsets:      [0, 128, 4224, 4736, 6784]

Executing FlashAttention directly over this contiguous buffer without unpack/re-pad overhead is the core objective of Jagged Flash Attention (JFA).

Hardware Evolution: Why NVIDIA Blackwell Demands TLX

The NVIDIA Blackwell architecture (B200) introduces fifth-generation Tensor Cores (TCGen05), enhanced Tensor Memory Accelerators (TMA), and asynchronous Warp-Group Matrix Multiply Accumulate (WGMMA) primitives. However, mapped hardware operations require strict memory alignment rules:

• TMA 2D/3D descriptors expect static strided memory layouts.
• Variable-length boundaries in jagged layouts trigger memory alignment violations if micro-tiles cross sequence boundary offsets.
• Warp-Group synchronization requires warp groups to cooperatively prefetch dynamic memory blocks into Shared Memory (SRAM) without causing thread divergence.

Writing pure CUDA or Triton kernels for jagged memory on Blackwell quickly leads to unmaintainable complexity. This is where Tile Language Extensions (TLX) enter the PyTorch infrastructure layer.

TLX provides Python and C++ abstractions that compile directly down to high-performance CUDA/SASS instructions. It allows kernel engineers to express layout transformations, TMA descriptor setups, and asynchronous double-buffering loops declaratively while granting fine-grained hardware layout control.

Architecture & Implementation of JFA via TLX

Jagged Flash Attention with TLX decouples sequence indexing from memory alignment. The kernel operates on micro-tiles (e.g., 128 x 64 or 64 x 128), dynamically calculating valid sequence boundaries using cu_seqlens while issuing hardware TMA read instructions.

Here is a conceptual implementation pattern showing how TLX expresses dynamic dynamic tile bounds and asynchronous memory execution for Jagged Flash Attention:

import torch
import torch.tlx as tlx

@tlx.jit
def jagged_flash_attention_kernel(
    Q_ptr, K_ptr, V_ptr, O_ptr,
    cu_seqlens_ptr,
    stride_q_tok, stride_k_tok, stride_v_tok,
    sm_scale: float,
    BLOCK_M: tlx.constexpr = 128,
    BLOCK_N: tlx.constexpr = 64,
    HEAD_DIM: tlx.constexpr = 128
):
    # Retrieve current batch index and block sequence index
    seq_idx = tlx.program_id(axis=0)
    head_idx = tlx.program_id(axis=1)

    # Load start and end boundaries for the specific sequence
    seq_start = tlx.load(cu_seqlens_ptr + seq_idx)
    seq_end = tlx.load(cu_seqlens_ptr + seq_idx + 1)
    seq_len = seq_end - seq_start

    # Determine dynamic tile ranges within bounds
    num_m_blocks = tlx.cdiv(seq_len, BLOCK_M)
    
    for m_tile_idx in range(num_m_blocks):
        # Calculate local offsets within the unpadded continuous buffer
        q_offset = (seq_start + m_tile_idx * BLOCK_M) * stride_q_tok
        
        # Load Query block using TMA descriptor abstraction
        q_tile = tlx.load_tma_async(
            Q_ptr + q_offset, 
            shape=[BLOCK_M, HEAD_DIM], 
            mask=(tlx.arange(0, BLOCK_M) < (seq_len - m_tile_idx * BLOCK_M))
        )
        
        # Initialize Softmax accumulators in Registers
        m_i = tlx.full([BLOCK_M], float("-inf"), dtype=tlx.float32)
        l_i = tlx.zeros([BLOCK_M], dtype=tlx.float32)
        acc = tlx.zeros([BLOCK_M, HEAD_DIM], dtype=tlx.float32)

        # Loop over Key and Value blocks for current sequence length
        num_n_blocks = tlx.cdiv(seq_len, BLOCK_N)
        for n_tile_idx in range(num_n_blocks):
            k_offset = (seq_start + n_tile_idx * BLOCK_N) * stride_k_tok
            v_offset = (seq_start + n_tile_idx * BLOCK_N) * stride_v_tok

            # Load K and V asynchronously into Shared Memory (SRAM)
            k_tile = tlx.load_tma_async(K_ptr + k_offset, shape=[BLOCK_N, HEAD_DIM])
            v_tile = tlx.load_tma_async(V_ptr + v_offset, shape=[BLOCK_N, HEAD_DIM])
            tlx.tma_wait()

            # Execute WGMMA (Warp Group Matrix Multiply Accumulate)
            qk = tlx.wgmma(q_tile, tlx.trans(k_tile)) * sm_scale
            
            # Online Softmax update logic
            m_ij = tlx.maximum(m_i, tlx.max(qk, axis=1))
            p = tlx.exp(qk - m_ij[:, None])
            l_ij = tlx.sum(p, axis=1)

            alpha = tlx.exp(m_i - m_ij)
            l_i = l_i * alpha + l_ij
            acc = acc * alpha[:, None] + tlx.wgmma(p.to(tlx.float16), v_tile)
            m_i = m_ij

        # Normalize and store final Attention output
        acc = acc / l_i[:, None]
        o_offset = (seq_start + m_tile_idx * BLOCK_M) * stride_q_tok
        tlx.store_tma_async(O_ptr + o_offset, acc.to(tlx.float16))

Key Technical Breakthroughs in TLX for Blackwell

1. Dynamic TMA Layout Descriptors

Traditional FlashAttention-2 and FlashAttention-3 rely on static strided layout assumptions where stride_batch = seq_len * num_heads * head_dim. In JFA, sequence lengths vary dynamically per batch element.

TLX handles this on Blackwell by generating runtime dynamic 2D TMA descriptors. By passing target memory descriptors directly to Blackwell's hardware address generation units, dynamic memory loads achieve native HBM3e bandwidth speed (>7.8 TB/s on B200) without triggering host-side kernel launch overheads.

2. Warp-Group Asynchronous Pipelining (WGMMA)

Blackwell splits execution into specialized producer and consumer warp groups:

• Producer Warps: Calculate dynamic segment bounds (cu_seqlens) and trigger asynchronous TMA loads into SRAM.
• Consumer Warps: Execute tensor core matrix operations (tcgen05.mma) directly on SRAM buffers.

Through TLX's pipeline semantics, memory transfer overhead is completely hidden behind double-buffered math execution.

Benchmarking: JFA on Blackwell (B200) vs. Legacy Implementations

To evaluate performance gains, tests were conducted comparing standard padded FlashAttention-2, native Triton JFA, and TLX-Optimized JFA on NVIDIA B200 GPUs (FP16 Precision, Head Dim = 128, Batch Size = 32, Average Seq Len = 2048 with 50% length variation variance).

Key Takeaway: TLX-optimized Jagged Flash Attention on Blackwell achieves nearly 1.88x speedup over naive JFA implementations and 2.5x speedup over standard padded implementations on B200, achieving up to 89% of theoretical memory bandwidth efficiency.

Practical Application & Developer Workflow

For AI application developers and enterprise platform managers, kernel-level optimizations directly translate to reduced API latency and lower per-token operational costs.

Developers accessing modern AI models via API gateways like n1n.ai benefit from these infrastructure innovations behind the scenes. When high-concurrency requests with variable prompt lengths hit inference clusters, underlying engines running optimized kernels like JFA ensure response times remain fast and steady.

Below is an example showing how PyTorch models natively call modern optimized attention APIs when configured with variable sequence lengths:

import torch
import torch.nn.functional as F

def run_jagged_attention_inference(q, k, v, cu_seqlens, max_seqlen):