FlashAttention

repository·main·Indexed 12 days ago

https://github.com/dao-ailab/flash-attention

An optimized Transformer implementation providing high-performance building blocks, including FlashAttention-4 (CuTeDSL) and fused CUDA kernels, to accelerate the training of models like GPT and ViT. It features optimized operations for matmul, bias, GELU, and LayerNorm, with specific support for bfloat16 and A100/Blackwell GPUs.

Tokens
11.9K
Snippets
31
Records
53
Agent score
98%

What's inside FlashAttention

  1. FlashAttention integration in ML frameworks

    main

    FlashAttention is integrated into several major machine learning frameworks and libraries, allowing for optimized Transformer operations without manual implementation. Key integrations include:

    • PyTorch: Integrated into the core nn.Transformer module.
    • Huggingface Transformers: Supported within the library.
    • Microsoft DeepSpeed: Integrated into the DeepSpeed inference engine.
    • NVIDIA Megatron-LM: Used for training large transformer language models at scale.
    • MosaicML Composer: Used for efficient neural network training.
    • EleutherAI GPT-NeoX: A research library for large language models.
    • PaddlePaddle: Accessible via the paddle.nn.functional.flash_attention API.
  2. Features of the LayerNorm CUDA extension

    main

    This CUDA extension provides a fused implementation of dropout, residual connections, and LayerNorm (building on Apex's FastLayerNorm). Key capabilities include:

    • Fused Operations: Combines dropout + residual + LayerNorm.
    • Architecture Support: Works for both pre-norm and post-norm architectures.
    • RMSNorm: Supports RMSNorm as an option.
    • Parallel Residuals: Supports layer norm with parallel residuals (e.g., used in GPT-J, GPT-NeoX, PaLM).
    • Dimension Support: Supports hidden dimensions up to 8192 (must be divisible by 8). For dimensions larger than 8k, contact the maintainers via an issue.
  3. Configure barrier tx_count for 2CTA / cluster mode

    main
    In 2CTA / cluster mode, multiple CTAs signal the same cluster-level mbarrier. If each CTA's TMA contributes N bytes, the expected transaction byte count (tx_count) must be N * cta_group_size (e.g., 2N for 2CTAs) rather than just N. This applies to all TMA pipelines (Q, K, and V).
  4. Choose the number of Warp Groups (num_wg) for SM90

    main

    The number of Warp Groups (num_wg) determines the tile_m size and the available register budget per thread. Increasing num_wg improves M-direction parallelism but increases shared memory (smem) usage and reduces the register budget.

    num_wgtile_m (fwd)ThreadsReg budgetBest for
    2128384216/threadhdim <= 128
    3192512128/threadhdim 129-192

    Hardware Constraints (H100):

    • SMEM: ~224 KB available for tensor buffers (after reserving ~3 KB for LSE, dPsum, and mbarriers).
    • GMMA atom: Always M=64. The effective M dimension (after swap) must be divisible by 64. The N dimension must be divisible by atom_layout_n * 8.
  5. Categorize debugging evidence by reliability tiers

    main

    Rank your evidence to understand the risk of corruption or hallucination:

    • Tier 1 (Trustworthy): Deterministic source-level facts, reproducible pass/fail results, or divergence against a reference implementation at a specific tensor index.
    • Tier 2 (Usable, needs corroboration): PTX (via CUTE_DSL_KEEP_PTX=1), dumped SASS (CUTE_CUBIN_PATH), shared-memory layout offsets, or cute.printf traces.
    • Tier 3 (Contaminated by definition): Anything captured at the moment of a trap (register values, faulting addresses, block/thread IDs, CUTE_DSL_LINEINFO attribution, or cuda-gdb backtraces).

    Warning: Tier 3 data is highly dangerous for building hypotheses because the corruption that caused the fault also produced the details that make the fault look like it fits your theory. Use Tier 3 only to confirm a prediction made using Tier 1 or 2 evidence.

  6. Distinguish between RS and SS modes in HGMMA

    main

    The second operand (Operand A) in an HGMMA instruction determines the memory access mode:

    • SS Mode (Shared Memory): Indicated by a shared memory descriptor like gdesc[UR..]. In this mode, both operands A and B are read from shared memory.
    • RS Mode (Register Source): Indicated by a plain register (e.g., R<N>). In this mode, operand A is read directly from registers, while operand B is read from shared memory.

    When to use RS: RS is used to reduce shared memory traffic when data is already present in registers from a previous computation (e.g., feeding dS from a softmax backward directly into a GEMM).

    # RS Mode: A=R232 (dS in regs), B=gdesc (K in smem)
    HGMMA.64x192x16.F32.BF16 R24, R232, gdesc[UR4], ...
    
    # SS Mode: both from smem
    HGMMA.64x192x16.F32.BF16 R24, gdesc[UR16], gdesc[UR4], ...
  7. Optimize dKV performance with mma_dkv_is_rs

    main

    The mma_dkv_is_rs (Register-Source for dKV) optimization can significantly reduce shared memory traffic and latency.

    Conditions for activation: When AtomLayoutMSdP == 1 AND AtomLayoutNdKV == num_wg AND SdP_swapAB is true AND dKV_swapAB is false.

    Benefits:

    • Eliminates sP from smem: Saves tile_m * tile_n * 2 bytes.
    • Eliminates P R2S store: Reduces shared memory traffic.
    • Eliminates A operand reads: Removes the need to read the A operand for dK and dV GEMMs by feeding P and dS matrices directly from registers.
  8. How padded_offset works for varlen buffers

    main

    In variable-length (varlen) operations, buffers such as dq_accum are laid out with tile-aligned gaps between sequences. The calculation of the padded_offset_q depends strictly on the tile_m value. If the tile size used for preprocessing differs from the tile_m used by the kernel, the preprocessing will zero out the wrong memory locations, leading to corrupted data in subsequent batches.

    padded_offset_q = ((offset_q + batch_idx * tile_m) // tile_m) * tile_m
  9. Understand AtomLayout and its impact on SMEM traffic

    main

    The atom_layout distributes Warp Groups (WGs) across the M and N dimensions of an MMA output.

    Given num_wg MMA WGs and atom_layout_m = A:

    • M direction: A warp groups, each handling M/A rows.
    • N direction: num_wg/A warp groups, each handling N/(num_wg/A) columns.

    Performance Trade-off:

    • Larger wg_n (more WGs in N direction): Each instruction reads a smaller B slice, but there are more total instructions reading overlapping A slices.
    • Smaller wg_n (fewer WGs in N direction): Fewer instructions total, but each instruction reads a larger B slice.
    • Rule of thumb: Smaller wg_n typically results in less total smem traffic.
  10. How R2P works in SM90 SASS masking

    main

    In SM90 (Blackwell) architecture, the R2P (Register to Predicate) instruction is used to optimize masking operations in FlashAttention kernels. Instead of using multiple ISETP (integer set-predicate) instructions to generate predicates for individual elements, a single R2P instruction can convert 7 bits from a register byte into 7 predicate registers simultaneously.

    The R2P Pattern

    The compiler generates a pattern where a bitmask is created via SHF, followed by R2P instructions to populate predicates, and FSEL instructions to apply those predicates for conditional masking (e.g., masking to -inf).

    Handling Bit Alignment

    Since R2P with the 0x7f immediate maps bits 0-6 to predicates P0-P6, the 8th bit (MSB) of each byte is not covered by the R2P instruction. For a 32-element row (4 bytes), the compiler handles the 4 leftover bits (bits 7, 15, 23, and 31) using separate LOP3.LUT or ISETP instructions to ensure full coverage.

    SHF.R.U32.HI R9, RZ, R9, R16    ; shift to create bitmask
    R2P PR, R9, 0x7f                  ; byte 0 → predicates P0-P6
    FSEL R15, R36, -INF, P6           ; apply P6: keep or mask to -inf
    R2P PR, R9.B1, 0x7f              ; byte 1 → predicates P0-P6
    FSEL R52, R52, -INF, P6           ; apply P6
    ... (repeated for bytes 2 and 3)
  11. Use swap_AB to handle non-divisible tile dimensions

    main

    The swap_AB option allows an MMA to swap its A and B operands, transposing the output tile. This is used to exchange which dimension maps to M (which must be divisible by 64) and which maps to N.

    When to use swap_AB:

    • If the natural M dimension is not divisible by 64, but N is (e.g., tile_m=80).
    • To change which operand resides in registers versus shared memory.

    Application in Backward Pass (5 MMAs):

    • SdP (S=Q@K^T, dP=dO@V^T): Swap if tile_m % 64 != 0.
    • dKV (dK=dS^T@Q, dV=P^T@dO): Swap if tile_n % 64 != 0 but hdim % 64 == 0.
    • dQ (dQ=dS@K): Swap if tile_m % 64 != 0 but hdim % 64 == 0.
  12. Identify GEMM dimensions from SASS instructions

    main

    You can derive the dimensions of a GEMM (General Matrix Multiply) by analyzing patterns in the SASS HGMMA instructions:

    1. K Dimension (Reduction): Count the number of consecutive instructions that use the same destination register. Multiply this count by 16 (the K-step size).
      • K = (count of instructions with same dest register) * 16
    2. M Dimension: Look for interleaved patterns of different destination registers. Each unique destination register represents a block of 64 rows (since M is fixed at 64 per instruction).
      • M = (number of unique destination registers) * 64
    3. N Dimension: This is explicitly provided in the instruction format (e.g., 64x96x16 implies N=96).
    # Example: dK = dS.T @ Q (192x96, K=64)
    # 3 accumulators (R120, R72, R24) -> M = 3 * 64 = 192
    # 4 instructions per accumulator -> K = 4 * 16 = 64
    
    HGMMA.64x96x16 dst=R120   # part 0, K-step 0
    HGMMA.64x96x16 dst=R72    # part 1, K-step 0
    HGMMA.64x96x16 dst=R24    # part 2, K-step 0
    HGMMA.64x96x16 dst=R120   # part 0, K-step 1
    HGMMA.64x96x16 dst=R72    # part 1, K-step 1
    HGMMA.64x96x16 dst=R24    # part 2, K-step 1
    ... (4 K-steps total)