M4.6 - FlashAttention-3/4 Hardware Pipelines: Hopper, Blackwell, and the Moving Bottleneck

0x00 Overview And Background

M4.3 derived how tiled online softmax removes the quadratic score and probability matrices from HBM. M4.4 then showed that an IO-aware algorithm can still leave the GPU underused because work partition, warp ownership, and non-matrix instructions matter. M4.5 moved to the asymmetric one-query decode shape.

This final article in M4’s execution/kernel branch asks a different question:

Once HBM traffic and coarse work partitioning have improved, how must attention be rescheduled around the asynchronous engines, on-chip memories, synchronization rules, and precision formats of a specific GPU generation?

The logical operation has not changed:

$$O = \operatorname{softmax}\left(\frac{QK^{\mathsf T}}{\sqrt{d_h}} + B + M\right)V.$$

Historical Problem: Every Improvement Exposes The Next Bottleneck

The FlashAttention lineage is not a sequence of interchangeable library versions. Each paper starts from a different residual bottleneck left visible by the previous algorithm and hardware schedule.

Starting point	Bottleneck that remains visible	Next design response	Detailed owner
Materialized attention	score and probability tensors move repeatedly through HBM	FA1 tiles exact attention and carries compact online-softmax state on chip	M4.3
FA1	IO improves, but non-matmul work, query-block parallelism, and warp ownership still limit utilization	FA2 reduces state-rescaling work and changes CTA/warp partitioning	M4.4
FA2 on Hopper	the schedule does not fully exploit TMA, WGMMA, warp specialization, and asynchronous overlap	FA3 overlaps movement, matrix work, and softmax and adds a hardware-aware FP8 path	this article, Hopper half
FA3-style Hopper schedule on Blackwell	tensor throughput and tile capacity grow faster than shared-memory and exponential throughput	FA4 moves accumulators into TMEM, redesigns forward/backward pipelines, and changes numerical scheduling	this article, Blackwell half

The FA3 paper motivates its Hopper redesign by reporting that FA2 reaches only about 35% utilization on H100 under the paper’s accounting convention. That is a source-reported coordinate, not a universal statement about every FA2 shape or implementation. The FA4 paper then identifies a different imbalance on Blackwell: tensor-core capability grows faster than shared-memory and special-function throughput, so a schedule tuned for Hopper exposes new bottlenecks.

The causal chain is:

remove quadratic HBM materialization
    -> expose work-partition and non-matmul overhead
    -> expose asynchronous-engine overlap and register pressure
    -> expose TMEM placement, shared-memory traffic, and exponential throughput

The execution machine has changed twice:

Hopper / FlashAttention-3:
    TMA moves tiles asynchronously
    WGMMA performs asynchronous warpgroup matrix operations
    accumulators mainly occupy registers
    producer and consumer warpgroups overlap movement, MMA, and softmax

Blackwell / FlashAttention-4:
    tcgen05 MMA writes accumulators into TMEM
    MMA tiles and tensor throughput grow
    exponential and shared-memory rates do not grow proportionally
    2-CTA MMA, TMEM allocation, software exponential work, and scheduling become central

The owner question is therefore not merely “which FlashAttention version is faster?” It is:

which hardware resource owns each intermediate?
which engine produces it?
which consumer is allowed to read it?
which dependency can overlap?
which resource becomes the next bottleneck?
what numerical contract changes in a low-precision path?
what evidence proves that the intended kernel ran?

Owner Contract: Owns, Consumes, And Defers

Owns

M4.6 derives generation-specific attention pipelines:

Hopper TMA, WGMMA, warpgroups, and register reallocation
FA3 producer/consumer and circular-buffer state
FA3 pingpong and two-stage GEMM-softmax overlap
FA3 FP8 layouts, block scales, and incoherent processing
Blackwell TMEM, tcgen05, and 2-CTA MMA
FA4 forward and backward resource-cycle derivations
partial exponential emulation and conditional online-softmax rescaling
persistent scheduling, deterministic reduction, and CuTe DSL boundaries
hardware/version-specific benchmark interpretation

Consumes

These are imported contracts, not invitations to repeat their full derivations:

Imported owner	Contract consumed here	New consequence derived in M4.6
M4.1	FLOP, byte, state, phase, and evidence-coordinate accounting	convert tensor work and source-reported rates into generation-specific resource ledgers
M4.2	CTA, warpgroup, occupancy, barrier, and profiler vocabulary	apply those contracts to TMA/WGMMA and TMEM/tcgen05 pipelines
M4.3	exact compact online-softmax state and merge algebra	assign the state to hardware roles and study delayed reference changes
M4.4	FA2 state coordinate, query-block ownership, warp partition, and backward baseline	identify the utilization pressure that FA3 and FA4 schedules address
M4.5	split-state reduction, cache/mask boundary, backend-admission schema, and profiler proof	record only architecture-specific artifact eligibility and selected-kernel evidence
M8.3	exact equivalent-transform theory	trace the attention-local transform before FP8 quantization
M8.6	live activation/KV scale lifecycle	bound the local FP8 scale and accumulator contract

Defers

Deferred owner	Material not re-derived here
M4.5	generic backend selection, one/few-query split reduction, paged address translation, and mask-program mechanics
M5	request scheduling, paged K/V allocation, placement, batching, and serving state
M8.2	model-wide calibration, sensitivity, and precision allocation
M8.3	general outlier geometry and equivalent-transform design
M8.6	live activation and K/V quantization, scale lifetimes, and state lifecycle
M8.9	quantized artifact compatibility, loader validation, and runtime admission
M8.10	deployment evaluation, promotion, release, and rollback of the exact quantized candidate
M4.7-M4.9	empirical training scaling, extrapolation, data constraints, downstream metrics, and lifecycle allocation
semantic long-context evaluation	whether a model can use distant evidence rather than merely process a longer tensor

The boundary can be summarized as:

M4.6 owns the physical schedule inside one generation-specific kernel.
It consumes earlier mathematical and execution contracts.
It does not own request policy, general backend policy, or model-wide quantization.

Claim Status And Source Discipline

Status	Meaning
Source-local	A19 material is preserved where it remains useful, especially the FA2-to-FA3 transition.
Externally verified	Mechanisms are checked against the FA3/FA4 papers and official NVIDIA Hopper, WGMMA, tcgen05, and CuTe documentation.
Paper algorithm	A mechanism or formula is scoped to the cited paper revision and its stated hardware.
Pinned implementation behavior	A source-code claim is scoped to an exact repository tag, commit, and path.
Current release behavior	Runtime support is scoped to a named release rather than inferred from the paper.
Rewrite derivation	Cycle, pipeline, register, and numerical identities are derived from source-reported rates and tensor shapes.
Source-reported	Throughput, speedup, accuracy, and compile-time results remain attached to the exact paper hardware and benchmark setup.
Boundary correction	A broad statement in A19 is narrowed when the primary source supports a more precise claim.

Evidence Epochs: Paper, Hardware, Source, Release

The same name can refer to a paper algorithm, an instruction contract, a source tree, or an installed runtime. M4.6 keeps those evidence surfaces separate.

Evidence surface	Immutable or dated coordinate used here	What it can establish	What it cannot establish
A19 local source	repository Markdown snapshot	the historical handoff and explanation style to recover	external correctness by itself
FA3 paper	arXiv 2407.08608v2	Hopper algorithm, FP8 method, and source-reported H100 results	current library support or a local speedup
FA4 paper	arXiv 2603.05451v1	Blackwell algorithm, resource model, and source-reported B200 results	every mode in the current artifact
NVIDIA documentation	official online documentation checked 2026-07-28	TMA, WGMMA, mbarrier, TMEM, tcgen05, CTA-pair, and DSMEM semantics	which FlashAttention specialization ran
Official implementation	fa4-v4.0.0.beta23, commit 2409214a03797b168f648ea30df1adbc09ce658a	version-scoped dispatch, role, pipeline, and feature behavior	timeless FA3/FA4 behavior
Local deployment	no GPU build, profile, or timing in this evidence pass	no performance claim	kernel selection, numerical behavior, or speed on this machine

Three rules follow:

paper algorithm != every later source path
source path exists != runtime selected it
source-reported benchmark != local deployment result

Corrections To The A19 Handoff

The short FA3 section in A19 is a useful historical pointer, but three claims need tighter boundaries:

not:
    Hopper is "up to 3x A100," therefore an attention kernel is 3x faster

instead:
    compare the measured kernel on a named GPU, precision, shape, and baseline

not:
    FA3 "uses QuIP"

instead:
    FA3 uses an orthogonal incoherent transform inspired by quantization work,
    including QuIP's line of ideas; this is one component of its FP8 path

not:
    a faster attention kernel creates semantic long-context ability

instead:
    it changes the runtime ceiling; model quality beyond the training regime
    is a separate position, data, and evaluation question

0x01 One Attention Equation, A Moving Hardware Bottleneck

FlashAttention-1 attacked HBM traffic. FlashAttention-2 improved work partitioning and reduced non-matrix overhead. That changes the denominator of the optimization problem: costs hidden behind HBM traffic become visible.

For one tiled forward iteration, the kernel repeatedly performs:

load K_j and V_j
S_j = Q_i @ K_j^T
P_j = online_softmax_update(S_j)
O_i = O_i + P_j @ V_j

A coarse resource model is:

$$T_{\text{tile}} \ge \max\left( T_{\text{load}}, T_{QK}, T_{\text{softmax}}, T_{PV}, T_{\text{store}} \right)$$

only when independent engines and dependencies permit full overlap. Without overlap, the same terms approach a sum:

$$T_{\text{serial}} \approx T_{\text{load}} + T_{QK} + T_{\text{softmax}} + T_{PV} + T_{\text{store}}.$$

This gap is the scheduling opportunity. It is also constrained by finite on-chip state. A deeper pipeline needs more shared-memory stages, more accumulator slots, more registers, more barriers, or more TMEM columns.

The Four Contracts Of A Hardware Pipeline

Contract	Question
Dependency	Has the producer completed the exact tile a consumer will read?
Lifetime	Can a buffer be reused, or is an older iteration still live?
Capacity	Do SMEM, registers, TMEM, and barrier slots fit at the chosen tile/stage count?
Numerical	Does reordered or lower-precision execution still implement the intended attention result within the stated tolerance?

An asynchronous instruction does not remove these contracts. It makes them explicit.

0x02 Hopper Prerequisites: TMA, WGMMA, And Warpgroups

Hopper provides two primitives that reshape the kernel schedule. This section does not repeat M4.2’s general GPU hierarchy. It specializes that hierarchy into the contracts an SM90 attention mainloop must actually satisfy.

The Hopper Resource Envelope

The official Hopper Tuning Guide gives the following compute-capability 9.0 limits for H100:

Resource	H100 limit used in this article	Why FA3 consumes it
register file	65,536 32-bit registers per SM	WGMMA accumulators, score fragments, online-softmax state, output state, pointers
architectural maximum	255 registers per thread	large consumer fragments approach the per-thread ceiling
resident thread blocks	at most 32 per SM	an architectural ceiling before registers and SMEM reduce residency
shared memory	228 KiB per SM	Q, staged K/V, optional P, barriers, and epilogue aliases
addressable shared memory	at most 227 KiB per block	CUDA reserves 1 KiB; allocations above 48 KiB require explicit dynamic-SMEM opt-in

These are capacity ceilings, not a residency result. A kernel still needs its compiled register count, rounded dynamic-SMEM allocation, thread count, cluster shape, and launch attributes. In particular, a legal 227 KiB block can leave only one such block resident, while a lower-SMEM specialization can still be register-limited.

Tensor Memory Accelerator

TMA performs asynchronous tensor transfers between global memory and shared memory. Official Hopper documentation describes one-dimensional through five-dimensional tensor copies, with address generation and bounds handling offloaded from ordinary per-element load instructions.

The important execution consequence is:

a small producer role can issue a large tile transfer;
consumer warps can spend their register and instruction budget on compute;
completion is communicated through pipeline barriers;

TMA does not mean data is immediately available after issue. The consumer still waits on the corresponding stage-complete state.

A Descriptor Is Not A Completed Tile

The host-side tensor map and the kernel-side TMA atom encode a multidimensional coordinate-to-address mapping. For one attention tile, the source-level transaction payload is:

$$X_Q = B_M d b, \qquad X_K = B_N d b, \qquad X_V = B_N d_v b,$$

where b is bytes per element. The pinned beta23 SM90 source computes these values with cute.size_in_bytes before constructing Q, K, and V TMA atoms. For B_M=B_N=d=d_v=128 and FP16/BF16, each transfer carries:

$$128\cdot128\cdot2 = 32{,}768\ \text{bytes} = 32\ \text{KiB}.$$

The transaction byte count is part of correctness, not only accounting. The PTX ISA TMA contract specifies a non-blocking cp.async.bulk.tensor operation whose global-to-shared completion can decrement the transaction count of an mbarrier. A consumer may treat the tile as complete only after the barrier phase observes both:

pending arrivals == 0
transaction count == 0

Calling arrive before the transfer finishes does not satisfy the second condition.

Full And Empty Are Different Barrier Facts

The pinned CuTe source allocates two barrier objects for Q and 2s each for K and V: one full and one empty state per stage. For s K/V stages:

$$N_{\text{mbarrier}} = 2 + 2s + 2s = 2 + 4s,$$ $$M_{\text{mbarrier,raw}} = 8(2+4s)\ \text{bytes}.$$

The bytes are small; the state machine is not. A full barrier proves that the producer’s transfer for (stage, phase) completed. The empty barrier proves that every required consumer released that same stage before it is reused. One cannot substitute for the other.

The circular pipeline coordinate is a stage index plus phase parity:

$$i'=(i+1)\bmod s,$$ $$\phi' = \phi \oplus \mathbf{1}[i=s-1].$$

Parity flips only when the index wraps. If a producer waits on the wrong parity, it can consume an old empty signal; if a consumer waits on the wrong parity, it can read a previous generation of the same SMEM address.

Shared Memory Is A Liveness Equation

At the pinned beta23 coordinate, the SM90 source builds one-stage Q storage, s-stage K/V storage, and optional shared P storage. Ignoring alignment and non-payload state:

$$|sQ|=B_Mdb, \qquad |sK|=sB_Ndb, \qquad |sV|=sB_Nd_vb,$$ $$|sP|=\mathbf{1}_{P\text{ in SMEM}}B_MB_Nb.$$

When Q is retained in registers, the source aliases the Q and staged-V arena after Q’s shared lifetime ends:

$$M_{\text{payload}}= \begin{cases} |sQ|+|sK|+|sV|+|sP|, & Q\text{ not retained in registers},\\ \max(|sQ|,|sV|)+|sK|+|sP|, & Q\text{ retained in registers}. \end{cases}$$

For the 128x128, d=d_v=128, FP16/BF16 example with s=2 and register-sourced P:

Q payload                 32 KiB
two K stages              64 KiB
two V stages              64 KiB
without Q/V alias        160 KiB
with Q/V alias           128 KiB
raw Q/K/V barriers         80 B

This is a source-level lower bound. Swizzled layout cosizes, 1024-byte alignment, optional shared P, scheduler state, and compiler/runtime allocation still have to be reconciled with the launched function.

Warpgroup Matrix Multiply-Accumulate

Hopper WGMMA is a warpgroup-wide asynchronous matrix operation. One warpgroup is four warps, or 128 threads. Official WGMMA documentation exposes asynchronous issue, commit groups, and wait groups. Operand B comes from shared memory; operand A can come from shared memory or registers; accumulators reside in registers.

thread:
    scalar execution agent

warp:
    32 threads executing together

warpgroup:
    4 warps / 128 threads cooperating on WGMMA

CTA:
    producer and consumer warpgroups plus shared pipeline state

The PTX WGMMA contract has four distinct steps:

1. establish operand visibility and register ordering
2. issue one or more wgmma.mma_async operations
3. collect prior uncommitted operations with wgmma.commit_group
4. use wgmma.wait_group N before consuming an accumulator whose group must finish

wait_group 0 drains all prior committed groups. wait_group 1 lets only the newest committed group remain pending and waits until every older one has completed, which creates overlap but also keeps another accumulator generation live. Reading a destination register before the group containing its WGMMA completes is undefined behavior.

Two fences solve different problems:

Ordering primitive	Orders	Does not prove
wgmma.fence	prior warpgroup register accesses before later WGMMA register accesses	that a TMA transfer completed
async-proxy fence	generic-proxy writes to a shared operand before the async proxy reads it	that WGMMA accumulator results are ready
wgmma.wait_group	completion of committed WGMMA groups	that another warpgroup reached a named scheduling rendezvous

The pinned source’s fence_view_async_shared() before a shared-memory P operand is therefore not interchangeable with warpgroup.wait_group(0). The first makes P visible to the async proxy; the second makes the WGMMA result safe to consume.

Register Reallocation Is A CTA Budget

Hopper also supports setmaxnreg, allowing producer warps to release register ownership and consumer warpgroups to request a larger maximum. The PTX instruction uses a per-CTA register pool; .inc can block until enough registers are available, and every warp in the warpgroup must execute the same operation. The immediate is an absolute per-thread maximum in [24,256], in multiples of eight.

At beta23, one producer warpgroup is paired with one, two, or three MMA warpgroups. The source requests:

MMA warpgroups	producer maximum	each consumer maximum	CTA ownership-hint sum
1	56	256	128*56 + 128*256 = 39,936
2	24	240	128*24 + 256*240 = 64,512
3	32	160	128*32 + 384*160 = 65,536

The last row exactly fills H100’s 65,536-register SM budget under this ownership arithmetic. It does not prove the kernel allocates those values, reaches a particular occupancy, or avoids spills. They are source-requested maxima. Compiled metadata remains required.

There is also a useful apparent mismatch to keep explicit: the Hopper guide lists an architectural maximum of 255 registers per thread, PTX permits a setmaxnreg immediate of 256, and the one-consumer source path requests 256. The safe interpretation is “source-level register-ownership hint at the legal PTX coordinate,” not “256 compiler-addressable registers were measured per thread.”

The reason a full 128-thread producer warpgroup enters the role branch even when one thread represents the TMA producer group is now visible: setmaxnreg is a warpgroup-collective contract. The beta23 code decreases registers for warps 0..3; inside the load routine, only the first producer warp handles the TMA path, and the TMA cooperative group is declared with size one.

Synchronization Is Plural

Domain	Source-level operation	Protected transition	Typical failure if omitted
TMA full/empty pipeline	producer_acquire, TMA issue, consumer_wait, consumer_release	GMEM-to-SMEM completion and stage reuse	partial tile read or early overwrite
WGMMA group	async GEMM plus wait_group	accumulator completion	consume an in-flight score/output fragment
async-proxy visibility	fence_view_async_shared	generic SMEM P write to WGMMA shared-operand read	stale P operand
compute-WG scheduler	named barrier arrive/sync	pingpong turn between consumer warpgroups	both groups enter an unsafe schedule phase
register redistribution	setmaxregister_decrease/increase	producer release and consumer acquisition of CTA register pool	undefined divergent execution or blocked increase

A generic sync in pseudocode hides too much. This section names the object and the dependency each synchronization protects.

Pinned Source-To-Instruction Trace

The current implementation coordinate is fa4-v4.0.0.beta23, commit 2409214a03797b168f648ea30df1adbc09ce658a. It is broader than the FA3 paper artifact, so the trace records branches rather than pretending there is one universal Hopper path.

Branches Before Instructions

The pinned flash_fwd_sm90.py establishes:

architecture family:
    SM90a family only

K/V movement:
    TMA by default
    non-TMA paged-K/V branch when requested

Q movement:
    TMA when the packed-GQA tile is representable
    cp.async-style fallback otherwise

P-V operand A:
    registers when mma_pv_is_rs
    shared memory otherwise

overlap:
    intra-warpgroup two-stage path or non-overlapped path

Therefore “the Hopper kernel uses TMA” is too broad. Q, K/V, and O each have their own admission decision, and paged or irregular layouts can change the movement path while the architecture remains SM90.

From Source Abstraction To Expected PTX Family

Layer in the pinned source	Source coordinate	Expected PTX family	Evidence limit
Q/K/V tensor-copy atoms and byte counts	flash_fwd_sm90.py lines 260-314	cp.async.bulk.tensor...mbarrier::complete_tx::bytes on admitted TMA paths	exact PTX/SASS variant requires compilation
full/empty pipelines with producer/consumer groups	flash_fwd_sm90.py lines 451-516	mbarrier init, arrive/expect-tx, parity wait, and release operations	CuTe/CUTLASS can change the concrete instruction sequence
producer and consumer register roles	flash_fwd_sm90.py lines 577-635	setmaxnreg.dec/inc.sync.aligned.u32	requested maxima are not compiled allocation evidence
QK/PV tiled MMA atoms	flash_fwd_sm90.py lines 96-118	wgmma.mma_async with source-selected register/shared operand modes	exact shape and instruction count depend on specialization
non-overlapped tile body	flash_fwd_sm90.py lines 1349-1408	QK WGMMA, wait, softmax, proxy fence when needed, PV WGMMA	source order is not final scheduled SASS order
intra-WG overlap body	flash_fwd_sm90.py lines 1411-1477	QK and previous PV groups in flight, then wait_group 1 and wait_group 0	achieved overlap requires a SASS/profiler check
pingpong scheduler barriers	flash_fwd_sm90.py lines 1480-1545	named barrier arrive/sync	does not itself prove balanced WG issue

The older pinned hopper/mainloop_fwd_sm90_tma_gmma_ws.hpp remains useful as the historical FA3 C++ mainloop: it names PipelineTmaAsync, transaction bytes, producer acquire/commit, consumer wait/release, and explicit warpgroup waits. The current CuTe Python path is a later implementation surface, not a line-for-line transcription of that file.

One K Stage, End To End

coordinate:
    producer state = (stage i, phase phi)

producer:
    wait until EMPTY(i, phi) is complete
    configure FULL(i, phi) to expect X_K bytes
    issue tensor copy gK_j -> sK[i]

TMA engine:
    transfer X_K bytes
    signal transaction completion on FULL(i, phi)

consumer:
    wait until FULL(i, phi) completes
    issue QK WGMMA using sK[i]
    commit WGMMA group
    wait until the QK group is safe to consume
    release EMPTY(i, phi) only after sK[i] has no remaining readers

advance:
    i = (i + 1) mod s
    flip phi when i wraps

The TMA barrier and WGMMA wait occur in the same tile trace but certify different events. This separation is the core hardware contract.

Proof Boundary

This trace proves that the pinned source contains architecture-specific abstractions and ordering points consistent with Hopper’s documented instruction families. It does not prove the exact PTX opcode variant, generated SASS order, compiled registers, spills, selected runtime specialization, or achieved overlap. Those require build artifacts and profiler evidence in later loops.

No local GPU artifact or timing in this evidence pass: every source-to-instruction edge above remains source-audited and unmeasured.

State Ownership Ledger

Object	Pinned SM90 source location	Producer	Consumer	Reuse condition
Q_i tile	one-stage SMEM, optionally retained in registers so its arena can alias V	TMA or Q fallback path	QK consumer warpgroups	Q consumers release its phase; aliased storage starts only after Q capture
K_j stage	s-stage circular SMEM	TMA or K/V fallback path	QK WGMMA	QK group no longer reads the stage, then consumer release
V_j stage	separate s-stage circular SMEM, possibly in the Q/V aliased arena	TMA or K/V fallback path	PV WGMMA	PV group no longer reads the stage, then consumer release
score tile S_j	FP32 consumer-register fragment	QK WGMMA	masks, modifiers, online softmax, P conversion	score-dependent scalar/vector work complete
probability tile P_j	register fragment or shared P arena	softmax/conversion path	PV WGMMA	PV has consumed the operand; proxy visibility required for SMEM mode
online max/sum	FP32 registers	softmax update	later iterations and epilogue	output tile complete
output accumulator O_i	FP32 consumer registers	PV WGMMA plus rescaling	next iteration/epilogue	final store preparation complete
full/empty barriers	2+4s raw mbarrier objects for Q/K/V	producer, TMA completion, consumers	opposite pipeline role	matching (stage,phase) completes

The table is the real pipeline. The instruction names are only its implementation.

0x03 FA3 Producer-Consumer Pipeline And Circular SMEM

The owner question is now narrower than the Hopper prerequisite question:

Given asynchronous movement and matrix engines, what state machine lets one warpgroup fill a finite shared-memory ring while other warpgroups consume it without reading early, overwriting live data, or deadlocking at the tail?

FlashAttention-3’s paper answers with a producer/consumer CTA and an s-stage circular SMEM buffer. The pinned beta23 implementation preserves that contract but specializes it further: Q has a one-stage pipeline, K and V have independent s-stage pipelines, their logical coordinates can be staggered by intra-warpgroup overlap, and the pipeline state survives across tiles assigned by the persistent scheduler.

Algorithm 1 Is A Contract, Not Literal Current Source

The paper’s Algorithm 1 gives the minimum dependency graph:

producer warpgroup:
    decrease its register ownership maximum
    load Q_i once
    for each logical KV tile j:
        acquire stage j mod s
        issue TMA(K_j, V_j)
        make the stage visible as full

consumer warpgroups:
    increase their register ownership maximum
    wait for Q_i
    for each logical KV tile j:
        wait for K_j
        complete Q_i K_j^T
        update online softmax state
        wait for V_j
        complete P_j V_j
        release stage j mod s

The current source should not be paraphrased as if every line above remained one operation:

Contract	FA3 Algorithm 1	Pinned beta23 CuTe path
Q lifetime	load once per query tile	independent one-stage Q pipeline with its own toggled phase
K/V storage	one conceptual paired stage	separate K-full/K-empty and V-full/V-empty pipelines at a shared stage coordinate
tile traversal	written as increasing j	dense path traverses admissible n_block values from right to left
load engine	TMA	TMA when admitted; paged fallback can use cooperative cp.async
consumer body	one QK-softmax-PV iteration	non-overlap body or split first-half/steady/last-half overlap body
work assignment	one CTA-level query tile	pipeline state persists through a persistent tile-scheduler loop

The paper is the algorithmic contract. The beta23 source is one release-specific realization of it.

One Logical KV Tile Has Several Coordinates

Let one scheduler work tile admit J K/V blocks. Define traversal ordinal r, not source block index, as:

$$r\in\{0,\ldots,J-1\}.$$

For the dense beta23 path:

$$n_r = n_{\max}-1-r,$$

so r=0 refers to the rightmost admissible block. Its circular-buffer coordinates are:

$$i_r = r\bmod s, \qquad e_r = \left\lfloor\frac{r}{s}\right\rfloor, \qquad \phi_r = e_r\bmod 2.$$

These coordinates answer different questions:

Coordinate	Question answered
work tile (m_block, head, batch)	which output tile owns this traversal?
n_r	which logical K/V block is being loaded?
optional page_idx(n_r)	which physical page supplies the block?
stage i_r	which SMEM slot holds it now?
epoch e_r	how many times has the ring wrapped?
parity phi_r	which generation of this stage is a barrier wait observing?

n_r and i_r are not interchangeable. A paged source can change the physical address while preserving the same stage coordinate; a later traversal can reuse the same stage while referring to a different logical block.

Source Encoding Of Index And Phase

The pinned PipelineStateSimple packs index and phase into one integer z:

$$i(z)=z\bmod s, \qquad e(z)=\left\lfloor\frac{z}{s}\right\rfloor, \qquad z' = z+1.$$

For reasoning, the barrier generation is the parity e(z) mod 2. The source returns the quotient e(z) directly and comments that an explicit modulo is omitted in this path. The evidence oracle normalizes the value to a parity bit rather than relying on that implementation shortcut.

Producer and consumer states deliberately start one generation apart:

consumer: z_c = 0  -> stage 0, full phase 0
producer: z_p = s  -> stage 0, empty phase 1

The opposite producer phase encodes that every ring entry is initially empty. For s=1, the source special-cases advancement by XORing the phase. For s>1, incrementing z changes stage each iteration and changes epoch only on wrap.

For s=2, J=5, and dense blocks [4,3,2,1,0]:

r	n_r	stage i_r	consumer raw phase	semantic full parity	producer raw phase	semantic empty parity
0	4	0	0	0	1	1
1	3	1	0	0	1	1
2	2	0	1	1	2	0
3	1	1	1	1	2	0
4	0	0	2	0	3	1

Stage zero appears three times, but each appearance carries a different generation. Waiting only on stage=0 would admit an ABA error: a consumer could mistake an old completion for the new payload.

Why A Circular Buffer Is Needed

Let P_X be the number of payloads committed to pipeline X and R_X the number released by its consumers. For each of the independent K and V rings:

$$0\le P_X-R_X\le s, \qquad X\in\{K,V\}.$$

The left inequality forbids releasing data that was never produced. The right inequality is the ring-capacity bound. Attempting to produce item r when r-s has not been released must block at producer_acquire; allowing it to proceed would overwrite a live generation.

At clean completion:

$$A_X=C_X=W_X=R_X=J,$$

where A, C, W, and R count producer acquire, producer commit, consumer wait, and consumer release events. Equality of counts is necessary but not sufficient: each event must also carry the same (stage, phase) and logical payload identity.

The producer can run ahead only while empty generations remain. The consumer can run only after the matching full generation completes. In the ideal steady state the two agents therefore exert backpressure on each other rather than using an unbounded queue.

Startup, Steady State, And Drain

For a two-resource model with per-tile load time T_L, per-tile dependent compute time T_C, at least two stages, and no synchronization overhead, the finite-J schedule is:

$$T_{\text{ideal}}(J) = T_L+T_C+(J-1)\max(T_L,T_C).$$

The derivation exposes the three phases:

fill:        first tile must load before its compute can begin
steady:      each later tile advances at max(T_L, T_C)
drain:       the final consumer work completes after the last load issue

The serial baseline is:

$$T_{\text{serial}} = J(T_{\text{load}} + T_{\text{compute}}),$$

so finite-work speedup is:

$$S_J= \frac{J(T_L+T_C)} {T_L+T_C+(J-1)\max(T_L,T_C)}.$$

Only as J grows does this approach:

$$S_{\infty} = \frac{T_{\text{load}} + T_{\text{compute}}} {\max(T_{\text{load}},T_{\text{compute}})}.$$

If T_L=300, T_C=500, and J=5, serial execution takes 4,000 idealized cycles while the two-stage lower bound is:

$$300+500+4\cdot500=2{,}800\ \text{cycles},$$

or 1.43x, not the asymptotic 1.6x. With J=1, there is no overlap benefit. With one SMEM stage, the next load cannot start before the current consumer releases the only slot, so this two-resource overlap model also collapses to the serial schedule.

These are dependency lower bounds, not H100 timings. TMA issue cost, WGMMA latency, barriers, masks, uneven K/V timing, occupancy, and scheduler work remain outside this model.

The Pinned Producer Trace

The beta23 producer initializes its K/V state once before the persistent scheduler loop. For each work tile it derives sequence and page coordinates, computes the admissible block interval, and walks that interval from right to left.

kv_write = producer_state(stages=s)       # source phase-index starts at s
q_phase = 1

while scheduler.has_work():
    work = scheduler.current()
    n_min, n_max = admissible_kv_interval(work)

    acquire(K, kv_write)
    load_and_commit(K[n_max - 1], stage=kv_write.index)
    acquire_and_load_Q(work, phase=q_phase)
    q_phase ^= 1

    acquire(V, kv_write)
    load_and_commit(V[n_max - 1], stage=kv_write.index)
    kv_write.advance()

    for n in reversed(range(n_min, n_max - 1)):
        acquire(K, kv_write)
        load_and_commit(K[n], stage=kv_write.index)
        acquire(V, kv_write)
        load_and_commit(V[n], stage=kv_write.index)
        kv_write.advance()

    scheduler.advance()

# Synchronization tail, not a dummy data tile.
producer_tail(V, kv_write)

That sketch is the non-overlapped TMA shape. The actual function also has paged cp.async, block-sparse, and staggered intra-warpgroup branches. In every branch, load_KV performs the source copy and then calls the matching pipeline’s producer commit.

The first s acquires still execute in source, but they should not block because the ring starts empty. This reconciles the paper’s statement that no waits are issued during fill with the pipeline abstraction’s unconditional-looking producer_acquire call: the logical wait returns immediately for an already-empty phase.

K And V Are Separate Reuse Proofs

The paper releases a conceptual paired K/V stage after PV. The pinned source can release the K generation earlier because QK is its final reader:

Payload	First legal consumer action	Last dependent action	Legal empty release
K at (i,phi)	full wait completes	QK WGMMA group completes	after wait_group makes the score fragment readable
V at (i,phi)	full wait completes	PV WGMMA group completes	after output accumulator update is complete

This separation permits the producer to reacquire a K slot while the corresponding V slot is still live. It does not permit advancing the logical K/V identity inconsistently. Both pipeline objects must agree on which traversal ordinal a stage generation represents.

The non-overlapped consumer body is therefore:

wait(K, read_state)
S = issue_QK(stage=read_state.index)
wait_wgmma_group(0)
release(K, read_state)

P, online_state = online_softmax(S, online_state)

wait(V, read_state)
O = issue_PV(P, V[read_state.index], O)
wait_wgmma_group(0)
release(V, read_state)
read_state.advance()

release(K) before QK completion is a use-after-reuse bug. release(V) before PV completion is the analogous output corruption bug. Advancing read_state does not release either barrier; advancement changes coordinates, while release changes ownership.

Intra-Warpgroup Overlap Splits The Loop Body

When intra_wg_overlap is enabled, the current source does not execute one complete QK-softmax-PV body per loop call. It peels the first and last halves:

first half:
    wait K_0 -> QK_0 -> release K_0 -> softmax_0

steady call at state r:
    clone state r for V_r
    advance state to r+1 for K_(r+1)
    issue QK_(r+1) and PV_r
    wait_group 1 -> release K_(r+1)
    softmax_(r+1)
    wait_group 0 -> release V_r

last half:
    wait V_last -> PV_last -> release V_last -> advance

This is why one shared kv_consumer_state can look off by one if read without the peeled-loop context. During the steady call, the cloned state names the previous V payload while the advanced state names the next K payload. 0x05 owns the performance and register-liveness consequences of that WGMMA-softmax overlap; this section owns the stage identities and release legality.

Tail Is Four Different Problems

“Handle the tail” is too vague for an asynchronous pipeline. At least four tails exist:

1. Ragged data tail. If sequence length is not divisible by tile_n, the rightmost logical block contains invalid score columns. The source isolates the first traversed block and applies sequence masking. TMA and non-TMA paged paths also have different address-safety behavior; a masked score does not by itself make an invalid page-table read safe.
2. Pipeline drain. After the last load, consumers must finish every produced K and V generation. No unproduced stage should receive a synthetic full signal merely to make the loop shape regular.
3. Producer exit tail. The pinned source calls pipeline_v.producer_tail(kv_producer_state) because V is the final produced dependency. CUTLASS defines producer tail as ensuring used buffers are synchronized before producer exit so dangling barrier signals do not survive. It is synchronization closure, not another V tile.
4. Persistent-work boundary. Q phase and K/V state are not reset for each scheduler tile. A work tile that produces no sparse blocks still releases Q and resets its numerical output state, but it must not invent K/V state advances.

Producer tail is not a payload. It changes no logical tile count and creates no consumer-visible data generation.

In other words, pipeline state persists across scheduler work tiles; the tile scheduler advances logical work without reinitializing the shared barrier generations.

For the dense branch, block interval, producer traversal, and consumer traversal must agree exactly:

$$J=n_{\max}-n_{\min}.$$

For block-sparse execution, J is the number of selected block coordinates, not the width of the dense interval. The conservation equality must use that selected count.

Barrier State Machine

EMPTY(i, producer_phase)
    -- producer_acquire returns --> OWNED_BY_PRODUCER
    -- issue async copy + commit --> IN_FLIGHT(i, consumer_phase)
    -- expected bytes complete --> FULL(i, consumer_phase)
    -- consumer_wait returns --> OWNED_BY_CONSUMER
    -- final dependent WGMMA completes --> RELEASABLE
    -- consumer_release --> EMPTY(i, next_producer_phase)

Full and empty are independent facts. WGMMA completion is a third fact. A full barrier proves that movement completed for one payload generation; it does not prove that a consumer has stopped reading it.

Deadlock Closure By A Decreasing Wait-For Graph

For one ring, producer item r can wait for release of r-s, while consumer item r can wait for producer commit of r:

produce(r) waits on release(r-s), only when r >= s
consume(r) waits on commit(r)
release(r) follows completion of consume(r)

Following a producer wait backward gives:

$$P_r \leftarrow R_{r-s} \leftarrow C_{r-s} \leftarrow P_{r-s}.$$

The ordinal decreases by s. The first s producer acquires have an initially empty base case, so this graph is acyclic if producer and consumer enumerate the same payloads and every completed consumer releases. This is the liveness proof behind the ring, not merely an implementation convention.

The proof fails under these bugs:

Bug	Broken invariant	Likely symptom
reset K/V state at every persistent work tile	phase continuity	wait on an old generation or overwrite a live stage
producer skips a selected block but consumer expects it	commit(r) exists	permanent consumer full wait
consumer exits a branch without V release	R_V=J	producer blocks when the ring wraps or in tail
release K before QK wait_group	last-reader completion	intermittent score corruption
release V before PV wait_group	last-reader completion	intermittent output corruption
advance without release	coordinate and ownership conflated	ring appears to progress but empty barrier never arrives
release a stage that was never produced	R_X<=P_X	later producer observes false availability
reuse stage index without generation parity	ABA protection	stale completion accepted as current data
omit producer tail in a clustered path	exit synchronization	dangling arrives or cluster-dependent hang

The dangerous failures are temporal rather than shape errors: every tensor can have the expected dimensions while containing the next tile, the previous tile, or a partially completed transfer.

The visual is a dependency schedule, not a measured timeline. The source-audited pipeline proves ordering and liveness obligations; compiled instruction overlap and barrier-stall cost still require a Hopper build and profile.

0x04 FA3 Pingpong: Overlap Softmax With Another Warpgroup’s GEMM

The producer-consumer analysis in 0x03 separated data readiness from data reuse. Pingpong adds another, independent order: which consumer warpgroup may issue its next tensor-core bundle. It does not move K/V payloads and it does not transfer online-softmax state between warpgroups.

Why A Small Operation Can Own The Schedule

Faster matrix units make softmax visible. The FA3 paper gives the following H100 SXM5 comparison for head dimension 128:

matmul work / exponential work:       about 512x
matmul throughput / special-function throughput: about 256x

Convert operation counts and rates into time:

$$\frac{T_{\exp}}{T_{\text{mma}}} = \frac{F_{\exp}/R_{\exp}}{F_{\text{mma}}/R_{\text{mma}}} = \frac{1}{512}\times 256 = \frac{1}{2}.$$

So even though there are far fewer exponentials, they can consume time equal to roughly half the matrix work in this source example. FP8 doubles the relevant tensor-core throughput while special-function throughput stays unchanged, making the ratio less favorable.

That ratio can be turned into an ideal scheduling bound. Let:

M = time for the ordered tensor-core bundle assigned to one consumer quantum
E = time for its scalar/MUFU softmax work

Two consumers executed serially take:

$$T_{\text{serial},2}=2(M+E).$$

If the second consumer’s tensor work overlaps the first consumer’s softmax perfectly, the two-item flow-shop lower bound is:

$$T_{\text{ping},2}=M+\max(M,E)+E,$$

and:

$$S_2= \frac{2(M+E)}{M+E+\max(M,E)}.$$

For a long stream of independent consumer quanta:

$$T_{\text{ping}}(N)=M+E+(N-1)\max(M,E),$$

so the asymptotic ceiling is:

$$S_{\infty}= \frac{M+E}{\max(M,E)} \le 2.$$

With the paper’s FP16 sketch, $E/M\approx 1/2$. The two-item ideal speedup is therefore 1.2x, while the long-stream ceiling is 1.5x. If FP8 halves $M$ while $E$ stays fixed, the two resources become closer to balanced and the asymptotic overlap opportunity approaches 2x. Neither number is a kernel prediction: barriers, instruction issue, dependencies, register pressure, tails, and imperfect compiler scheduling all reduce it.

The paper reports movement from roughly 570 TFLOP/s to 620-640 TFLOP/s for one FP16 forward configuration (d_h=128, sequence length 8192), or about 1.09-1.12x. It also reports 661 TFLOP/s when its tested optimizations are combined in another fixed configuration. These are source-reported ablations, not a universal pingpong multiplier.

Pingpong Orders Issue; It Does Not Exchange Numerical State

The paper’s picture can look like two warpgroups alternately compute the same output. That is not the useful implementation model. Each consumer warpgroup owns its own partitioned fragments and preserves its own online-softmax recurrence:

State or event	Owner	What crosses warpgroup boundaries?
local Q/S/P fragments	one consumer warpgroup	nothing
running row maximum, row sum, and output accumulator	one consumer warpgroup	nothing
K/V stage readiness and reuse	Q/K/V mbarrier pipelines from 0x03	full/empty generation signals
WGMMA completion	issuing warpgroup	WGMMA commit/wait ordering only
pingpong permission	one named-barrier token	issue-order credit, not tensor data

The current source partitions the tiled QK and PV operations by a warp_group_thread_layout. A consumer group advances the attention state for its own query-row partition. The named scheduler only delays the next tensor bundle long enough to encourage this pattern:

consumer WG0:
    issue GEMM bundle -> perform local softmax

consumer WG1:
    wait for token -> issue GEMM bundle while WG0 uses scalar/MUFU resources

then pass the issue token back

The phrase “roles swap” means the hardware-resource overlap swaps. It does not mean WG0 hands WG1 its $(m,\ell,O)$ state.

Paper, Historical Kernel, And Current Beta Are Three Evidence Surfaces

Surface	Scheduler contract	Boundary
FA3 paper Figure 1	two consumer warpgroups; bar.sync encourages one group’s softmax to overlap the other’s GEMMs	conceptual schedule and source-reported ablation
historical Hopper C++ mainloop at the pinned commit	two- or three-group named-barrier ring; explicit first signal; non-overlap and intra-WG branches	source behavior, not compiled proof
beta23 CuTe flash_fwd_sm90.py	one to three MMA warpgroups; scheduler barrier enabled only for selected compile-time cases	current pinned implementation behavior

The current beta computes:

use_scheduler_barrier = (
    num_wg_mma >= 2 and tile_hdim <= 128
    if intra_wg_overlap
    else num_wg_mma == 2
)

That gives the following admission table:

Consumer warpgroups $G$	Intra-WG overlap	tile_hdim <= 128	Named scheduler
1	either	either	disabled
2	false	either	enabled
3	false	either	disabled in this branch
2 or 3	true	yes	enabled
2 or 3	true	no	disabled

This is why “FA3 uses pingpong” is too broad. The source has a compile-time scheduler decision, and the selected branch depends on warpgroup count, head dimension, and whether intra-warpgroup overlap is enabled.

Physical Warpgroup, Local Group, And Barrier ID

The CTA uses one physical producer warpgroup before the consumer groups. Define:

$$g_{\text{phys}}\in\{1,\ldots,G\}, \qquad g=g_{\text{phys}}-1\in\{0,\ldots,G-1\}.$$

The beta23 forward enum reserves:

barrier 0: reserved for CTA-wide synchronization
barrier 1: Epilogue
barrier 2: WarpSchedulerWG1
barrier 3: WarpSchedulerWG2
barrier 4: WarpSchedulerWG3
barrier 5: PFull
barrier 6: PEmpty

Let $B_0=2$ be the first scheduler barrier. Consumer $g$ waits on:

$$B_g=B_0+g.$$

The source expression is the same mapping in physical coordinates:

$$B_g=B_0-1+g_{\text{phys}}.$$

Each consumer warpgroup has $W=128$ threads. A scheduler generation pairs the current group with one signaling group, so the named-barrier participant count is:

$$N_{\text{barrier}}=2W=256.$$

This count is not interchangeable with the Q/K/V mbarrier arrival count. NVIDIA’s CUTLASS documentation explicitly distinguishes them: all participants in a named-barrier synchronization count, while an mbarrier pipeline can count only its signaling agents.

A Named Barrier Is A Reusable Generation, Not A Boolean Flag

PTX exposes 16 CTA named barriers, IDs 0..15. barrier.sync contributes the executing threads and waits; barrier.arrive contributes the executing threads without waiting. When the expected count is reached, the barrier releases waiters and automatically reinitializes for its next generation.

For scheduler barrier $B_g$, a normal generation combines:

128 threads from consumer g executing sync(B_g, 256)
128 threads from predecessor pred(g) executing arrive(B_g, 256)

where:

$$\operatorname{pred}(g)=(g-1)\bmod G, \qquad \operatorname{next}(g)=(g+1)\bmod G.$$

The first generation of $B_0$ has no predecessor yet. mma_init() therefore deposits one bootstrap half-credit on $B_0$ from the first consumer group. Its first sync(B_0,256) supplies the other 128 arrivals and starts the ring.

Conceptually, if $h$ is a scheduler handoff ordinal:

$$\begin{aligned} B_0^{(0)} &: \operatorname{init\_arrive}(0)+\operatorname{sync}(0,0),\\ B_g^{(h)} &: \operatorname{arrive}(g-1,h)+\operatorname{sync}(g,h), && g>0,\\ B_0^{(h+1)} &: \operatorname{arrive}(G-1,h)+\operatorname{sync}(0,h+1). \end{aligned}$$

At a stable persistent-work boundary, one half-credit remains pending on $B_0$. All other scheduler barriers are empty. The scheduler token is therefore conserved:

$$C_{\text{pending barrier credit}}+C_{\text{warpgroup owner}}=1.$$

While one group owns the token between sync completion and its arrive(next), no pending barrier credit exists. After arrive(next), ownership is encoded as one pending half-credit on the next group’s barrier.

Pinned Source Trace: Bootstrap, Wait, Issue, Pass

The current source reduces to this scheduler skeleton:

def scheduler_init(physical_wg):
    if use_scheduler_barrier and physical_wg == 1:
        barrier_arrive(B0, number_of_threads=256)

def wait_for_turn(physical_wg):
    local_wg = physical_wg - 1
    barrier_sync(B0 + local_wg, number_of_threads=256)

def pass_turn(physical_wg, G):
    local_wg = physical_wg - 1
    next_wg = (local_wg + 1) % G
    barrier_arrive(B0 + next_wg, number_of_threads=256)

For two consumer groups, one circuit is:

Order	Event	Barrier state after event
0	WG0 bootstrap arrive(B0)	half-credit pending on B0
1	WG0 sync(B0) completes	WG0 owns token
2	WG0 issues its tensor bundle	WG1 still blocked on B1
3	WG0 arrive(B1)	half-credit pending on B1
4	WG1 sync(B1) completes	WG1 owns token
5	WG1 issues its tensor bundle	WG0 can perform non-tensor work
6	WG1 arrive(B0)	token rests on B0 for the next circuit

For three consumer groups, the same source becomes B0 -> B1 -> B2 -> B0. Figure 1 remains a two-group explanation; it should not be used as proof that every current source specialization launches exactly two consumers.

The Non-Overlap Path Has Startup, Middle, And Drain Grants

Without intra-warpgroup overlap, the source places one scheduler sync before the first block, one arrive after QK, another sync before PV, and a final arrive after the last PV:

wait_for_turn()                       # startup grant for QK(0)

for j in logical_kv_blocks:
    issue_QK(j)
    pass_turn()                       # peer may issue its tensor bundle

    finish_QK_and_run_softmax(j)
    wait_for_turn()                   # returned grant
    issue_and_finish_PV(j)
    # The next loop immediately issues QK(j+1) before passing again.

pass_turn()                           # drain grant after PV(last)

For $J\ge1$ K/V blocks, each consumer executes:

$$H_{\text{non-overlap}}(J)=J+1$$

named-barrier syncs and the same number of arrives. The grants cover:

grant 0:       QK(0)
grant 1..J-1: PV(j-1) followed by QK(j)
grant J:       PV(J-1), then pass the drain token

The final arrive is not optional cleanup. Without it, the peer can remain blocked in its final sync or the next persistent work tile can start with the token stranded in the wrong group.

The Intra-Warpgroup Path Moves The Barriers Into The Steady Pair

The peeled overlap path from 0x03 has a different scheduler count:

startup:
    QK(0) outside the scheduler ring

steady j = 1..J-1:
    wait_for_turn()
    issue QK(j) and PV(j-1)
    pass_turn()
    finish dependencies and update softmax

drain:
    PV(J-1) outside the scheduler ring

The scheduler handoffs per consumer are therefore:

$$H_{\text{intra-overlap}}(J)=\max(J-1,0).$$

For J=1, the branch performs no steady-state scheduler handoff. The bootstrap half-credit remains on B0, ready for a later work tile. This is legal precisely because initialization is CTA-scoped rather than work-tile-scoped.

This subsection only owns where the cross-warpgroup token is consumed and passed. The register-buffer dependency that makes QK(j) and PV(j-1) legal remains the 0x05 owner.

Persistent Work And Sparse Branches Must Preserve The Token

In the pinned source, mma_init() executes once before the persistent while work_tile.is_valid_tile loop. Scheduler barriers are not reinitialized for every output tile. A full token-ring circuit returns one pending half-credit to B0, which the next work tile consumes when it first enters a scheduled steady region.

Dense, causal, local, and block-sparse branches can execute different numbers of arithmetic operations, but every participating consumer warpgroup must make the same scheduler decision and execute the same number of scheduler calls for a work tile. The block-sparse helper makes this explicit:

non-empty sparse work, non-overlap:
    one initial sync
    matched per-block scheduler calls
    one final arrive

empty sparse work:
    every consumer skips the scheduler calls together
    pending token remains unchanged

intra-overlap sparse work:
    first/steady/last split preserves the steady handoff count

The historical Hopper mainloop contains an even sharper warning. An append-K/V detour that temporarily reuses the first scheduler barrier consumes the pre-seeded WG0 credit before the detour and re-signals it afterward. The bootstrap arrival is state, not a harmless annotation.

Cross-Warpgroup Liveness Proof

Let the forward scheduler rank be:

$$\rho(h,g)=hG+g.$$

For $g>0$, consumer $(h,g)$ can leave its sync only after $(h,g-1)$ passes the token. Consumer $(h+1,0)$ depends on $(h,G-1)$. Every non-bootstrap wait therefore points to the immediately preceding rank:

$$\operatorname{start}(h,g) \leftarrow \operatorname{arrive}(\rho(h,g)-1) \leftarrow \operatorname{issue}(\rho(h,g)-1) \leftarrow \operatorname{start}(\rho(h,g)-1).$$

The base case is the bootstrap half-credit for $\rho(0,0)=0$. Because each recursive dependency decreases $\rho$ by one, the wait-for graph is acyclic under these obligations:

all 128 threads in a participating warpgroup execute each barrier call convergently
every granted group eventually reaches exactly one pass_turn for that grant
all groups choose the same scheduler-enabled branch and work-count path
the barrier ID and participant count are identical for both halves of a generation
WGMMA and K/V dependencies needed before pass_turn eventually complete

The proof establishes protocol liveness, not scheduler fairness or performance. A compiler can still reduce overlap, a long softmax can still leave tensor units idle, and a resource stall can delay the group that owns the token.

Why One Named Barrier Per Consumer Group

Design	Benefit	Cost or failure pressure
no scheduler barrier	no barrier instruction overhead	tensor bundles can cluster and leave MUFU/tensor overlap to incidental scheduling
CTA-wide barrier	simple global phase	includes unrelated producer work and destroys fine-grained overlap
one named barrier reused for every group	fewer IDs	generations are ambiguous when a fast group reaches the next turn early
per-group named-barrier ring	explicit predecessor/successor and automatic generation reset	consumes $G$ of 16 hardware barrier IDs and requires exact call balance
three-group ring	more independent consumer work	lower per-group register budget and more scheduling/resource pressure; not automatically faster
intra-WG overlap only	fewer cross-group handoffs	loses deliberate GEMM/MUFU ordering between consumers
combined intra- and inter-WG scheduling	exposes both dependency dimensions	most complex startup, steady, drain, and proof surface

An mbarrier array could encode a more general scheduler, but it would introduce shared-memory state, explicit phase handling, and different participant accounting. The pinned kernel uses mbarriers for payload readiness and named barriers for issue order; preserving that separation makes the proof auditable.

Scheduler Failure Matrix

Failure	Immediate effect	Later symptom
omit bootstrap arrival	every consumer waits on an empty ring	CTA deadlock at first scheduled bundle
use 128 instead of 256 participants	a group can complete its own barrier generation	ordering collapses; peer arrival enters the wrong generation
use more than 256 participants	the two consumer halves cannot satisfy the count	permanent named-barrier wait
map physical WG directly without subtracting producer WG	wait and arrive target different IDs	one barrier starves while another accumulates arrivals
one group skips a conditional barrier	generation counts diverge	shape- or mask-dependent deadlock
omit non-overlap final arrive	token remains owned after final PV	peer drain or next persistent tile hangs
re-prime every work tile	duplicate half-credit enters an active generation	phase drift, premature release, or deadlock
treat named barrier as K/V readiness	issue order is mistaken for payload completion	stale shared-memory read or early stage release
assume the paper timeline is compiled fact	performance claim lacks artifact proof	correct kernel runs slower than the clean diagram predicts

Runtime Proof Surface

The source predicts the protocol, but a Hopper artifact is needed to prove it survived compilation. A complete runtime record should include:

compile-time branch:
    num_wg_mma, intra_wg_overlap, tile_hdim, use_scheduler_barrier

artifact:
    BAR.SYNC / BAR.ARRIVE family, barrier IDs, participant count
    WGMMA issue/commit/wait placement around those barriers

dynamic balance:
    work tiles and logical K/V blocks per CTA
    expected scheduler handoffs for the selected branch
    barrier-stall and issue-stall observations

resource overlap:
    tensor-pipe activity
    MUFU/special-function activity
    timeline evidence that the activities overlap

correctness:
    dense, causal, local, sparse, empty, and ragged cases
    repeated persistent work tiles

For the non-overlap branch, the source-level expected count per consumer and non-empty work tile is J+1 syncs and J+1 arrives. For the intra-overlap branch it is max(J-1, 0) of each. These are static source expectations, not measured instruction counts; compiler specialization and disabled scheduler branches must be recorded before comparison.

Source: Shah et al., FlashAttention-3, Figure 1. The diagram is an intended schedule, not a guarantee that every issued instruction overlaps perfectly after compilation.

The reconstruction makes the source protocol visible; it is not a measured instruction timeline. The original paper figure explains the intended resource overlap, while the reconstruction explains how the pinned source tries to enforce progress.

0x05 FA3 Two-Stage WGMMA-Softmax Pipeline

Pingpong in 0x04 orders two consumer warpgroups. This section asks a different question inside one consumer warpgroup:

How can the kernel overlap WGMMA with online softmax without reading an unfinished accumulator, overwriting an in-flight register operand, or spending so many registers that the schedule loses residency?

The answer is not merely “add another stage.” A legal two-stage schedule needs four contracts at once:

algorithm:
    QK(j) -> softmax(j) -> PV(j)

WGMMA queue:
    commit QK before PV; wait by group age

register liveness:
    S_next, P_current, and O coexist across the overlap window

compiler evidence:
    the intended source order must survive lowering without spills or serialization

Paper Algorithm, Current Source, And Compiled Artifact

FA3 Algorithm 2 is the paper contract. It peels the first QK/softmax and the final PV, then overlaps the middle pairs. The current fa4-v4.0.0.beta23 CuTe source implements the same dependency shape with different abstractions: mma_qk_fn, mma_pv_fn, PipelineAsync K/V states, and explicit warpgroup.wait_group calls.

Those are related but not interchangeable evidence surfaces:

Surface	What it establishes	What it does not establish
FA3 Algorithm 2	intended two-stage numerical and wait order	behavior of every later implementation
beta23 source at commit 2409214...	one concrete source path and compile-time alternatives	emitted PTX/SASS or selected runtime kernel
resolved CuTe/Quack/compiler environment	exact helper and lowering versions	runtime overlap by itself
PTX/SASS artifact	actual instruction families, order, registers, and spills	dynamic overlap or end-to-end speedup
profiler trace	execution and stalls for one kernel coordinate	universal behavior across shapes

This section is paper-verified and source-audited, but locally uncompiled and unmeasured.

Derive The Two-Stage Schedule From The Dependency DAG

For one query block i and K/V traversal tile j:

$$S_j = Q_iK_j^{\mathsf T}.$$ $$\begin{aligned} (P_j,m_j,\ell_j) &=\operatorname{OnlineSoftmax}\!\bigl( \\[-2pt] &\qquad S_j;\,m_{j-1},\ell_{j-1} \bigr). \end{aligned}$$

Writing the rescaled old output accumulator as $R_j$ makes the last dependency two explicit steps:

$$R_j = \alpha_j O_{j-1}.$$ $$O_j = R_j + P_jV_j.$$

The three operations are not mutually reorderable:

QK(j) must finish before S_j is read by softmax.
softmax(j) must produce P_j before PV(j) reads P_j.
PV(j) must finish before O is rescaled, reused, or stored.

The legal overlap comes from adjacent tiles. After peeling tile 0, steady pair j performs:

older numerical state:    P_(j-1), V_(j-1), O
newer numerical state:    Q, K_j, S_j

1. issue and commit QK(j)
2. issue and commit PV(j-1)
3. wait until QK(j) is complete, while PV(j-1) may remain pending
4. compute softmax(j) from S_j
5. wait until PV(j-1) is complete
6. convert/store P_j and update the online-softmax output scale

The complete finite schedule is therefore:

Phase	Matrix work	Scalar/register work	Boundary
Startup	QK(0)	softmax(0), form P_0	no older PV exists
Steady j=1..J-1	QK(j) then PV(j-1)	softmax(j) overlaps the pending PV	two WGMMA groups
Drain	PV(J-1)	final rescale/store	no newer QK exists

Let M_Q, M_V, and E denote normalized time for QK WGMMA, PV WGMMA, and the scalar softmax/update path. A serial traversal has the dependency lower bound

$$T_{\text{serial}}(J)=J(M_Q+E+M_V).$$

Because softmax cannot begin until QK completes, the steady pair is bounded by QK followed by the slower of softmax and PV:

$$T_{\text{steady pair}}\ge M_Q+\max(E,M_V).$$

Including startup and drain gives the idealized finite bound

$$T_{\text{2-stage}}(J) \ge M_Q+E+(J-1)\left[M_Q+\max(E,M_V)\right]+M_V.$$

For M_Q=M_V=E=1 and J=5, serial work is 15 normalized units while the dependency-only two-stage bound is 11, a 15/11 ~= 1.36x finite speedup. This is a scheduling bound, not a benchmark prediction: issue bandwidth, WGMMA latency, barriers, occupancy, and compiler placement are omitted.

Source: Shah et al., FlashAttention-3, Figure 2.

wait_group(1) Is An Age Proof, Not A QK Label

WGMMA groups are ordered commit groups; they are not named “QK” and “PV” by the hardware. The pinned helper performs:

warpgroup.fence()
cute.gemm(...)
warpgroup.commit_group()
if wg_wait >= 0:
    warpgroup.wait_group(wg_wait)

The overlap path calls both GEMMs with wg_wait=-1, so the helper commits each group but leaves waiting to the caller:

oldest                                              newest
group g:     QK(j)  committed first
group g+1:   PV(j-1) committed second

wait_group(1): at most one newest group may remain pending
               -> QK(j) is complete
               -> PV(j-1) may still be pending

wait_group(0): no prior groups may remain pending
               -> PV(j-1) is complete too

The official PTX contract makes two uses undefined before a wait that covers their WGMMA group: accessing the accumulator, and accessing or reusing a register fragment used as matrix A. Applied to this schedule:

Resource/action	Earliest legal point	Reason
read or transform S_j	after wait_group(1)	QK is the older completed group
release shared K_j stage	after wait_group(1)	QK has finished reading K
read/rescale O after PV issue	after wait_group(0)	PV accumulates into O asynchronously
overwrite register P_(j-1) on the RS PV path	after wait_group(0)	P is PV’s register-A fragment
release shared V_(j-1) stage	after wait_group(0)	PV has finished reading V

wgmma.fence solves a different problem: it orders earlier register accesses against a later WGMMA access to the same accumulator or register-A fragment. Shared-memory producers additionally need the async-proxy visibility protocol derived in 0x02-0x03. A fence is not completion; a wait is not a proxy fence.

Pinned Beta23 Trace: Two Commits, Two Different Waits

The current source’s mma_one_n_block_intrawg_overlap path has this exact shape:

Source step	Group state	Live-state consequence
wait for next K; scheduler sync	no new WGMMA group	K stage becomes readable
mma_qk_fn(..., wg_wait=-1)	commit QK as older group	acc_S exists but is pending
optional rescale_O_before_gemm	only QK pending	O is still legal to touch
wait for previous V	QK pending	V stage becomes readable
mma_pv_fn(..., wg_wait=-1)	commit PV as newer group	O and P become in-flight operands
warpgroup.wait_group(1)	QK complete; PV may remain	release K; consume acc_S
score modifiers, mask, online softmax	PV may remain	scalar path overlaps the newer group
warpgroup.wait_group(0)	all prior WGMMA complete	release V; reuse P; touch O
convert acc_S -> tOrP	no prior PV uses tOrP	prepare probability operand for next pair

This explains an otherwise surprising source detail on the register-shared (RS) PV path: conversion into the persistent tOrP register fragment occurs after wait_group(0), not immediately after softmax. Reusing that fragment while the previous PV group still reads it would violate the PTX register-input contract. On the shared-memory (SS) PV alternative, converted P is copied through a shared-memory path with its own proxy-fence and visibility obligations; the persistent register-A interval below applies specifically to RS PV.

The beta package does not completely pin this helper implementation. Its dependency is quack-kernels>=0.5.3, without an upper bound. The audited 0.5.3 wheel implements wg_wait=-1 as “commit with no helper wait,” and the same core wrapper is present in 0.6.1; a reproducible compiled artifact must still record the actually resolved Quack wheel and hash.

Register Liveness Is The Capacity Contract

The important question is not whether a tensor is conceptually temporary. It is whether its live interval intersects the overlap window.

resource             live interval in a steady pair
-------------------  -------------------------------------------------
O accumulator        whole query traversal; in-flight during PV
P_(j-1) register A   before PV issue through wait_group(0)
S_j accumulator      QK issue through softmax and P_j conversion
row m and l          whole query traversal
row_scale            softmax(j) through the selected O-rescale point
K_j shared stage     K wait through wait_group(1), then releasable
V_(j-1) shared stage V wait through wait_group(0), then releasable

Let G consumer warpgroups cooperate on a row-partitioned tile, so the average ownership denominator is 128G consumer threads. For the RS PV path, a logical 32-bit-register-word lower bound is:

$$R_O^{\text{logical}}=\frac{B_r d_v}{128G}, \qquad R_{S}^{\text{logical}}=\frac{B_rB_c}{128G}, \qquad R_P^{\text{logical}}=\frac{B_rB_c b_P}{4\cdot128G},$$

where b_P is bytes per probability element. These are payload averages, not compiler register counts. Fragment layouts, padding, predicates, descriptors, row reductions, addresses, and compiler temporaries add state; uneven ownership can change the per-thread maximum.

For B_r=B_c=d_v=128, G=2, FP32 O/S, and FP16 P:

Live payload	CTA-wide bytes	Average 32-bit words per consumer thread
O	64 KiB	64
S_j	64 KiB	64
P_(j-1)	32 KiB	32
RS overlap subtotal	160 KiB	160

That 160 is a logical lower bound for the central RS live payload, not a prediction that ptxas will report exactly 160 registers. An SS PV path moves the persistent P operand into shared memory but adds store, visibility, and shared-memory-capacity costs. The pinned source’s two-consumer budget hint is 240 registers per consumer thread; for RS PV, the difference must cover online-softmax state, descriptors, masks, addresses, fragment-layout overhead, and compiler temporaries. A spill or one extra live fragment can erase the expected overlap benefit.

The paper describes the two-stage increment through the extra score accumulator:

$$M_{S_{\text{next}}} = B_rB_c\operatorname{sizeof}(\text{float}).$$

For B_r=B_c=128 and FP32 score accumulators:

$$128\times128\times4 = 65{,}536\ \text{bytes} = 64\ \text{KiB}.$$

That is 16,384 logical 32-bit values across the CTA, or 64 values per consumer thread under the G=2 even-ownership coordinate. The actual maximum remains a compiled property.

The reconstruction combines the paper algorithm, official PTX rules, and the pinned beta23 source. It is a dependency and liveness proof, not a captured GPU timeline.

Rescale Before Or After PV Is A Real Design Choice

The online update needs O expressed in the new row-max reference frame. The pinned source has two legal placements:

rescale_O_before_gemm = true:
    rescale O after QK issue but before PV issue
    keep scores_scale row state for the next update

rescale_O_before_gemm = false:
    issue PV first
    wait_group(0)
    rescale the now-complete O with row_scale

The first placement overlaps O rescaling with QK but must finish before O becomes PV’s in-flight accumulator. The second shortens pre-PV scalar work but puts rescaling after the full drain. They are algebraically coordinated with the source’s scores_scale state; moving the instruction without moving the reference-frame bookkeeping is a numerical bug, not a harmless schedule change.

Why The Paper Stops At Two Stages

A three-stage proposal overlaps:

QK(j+2), softmax(j+1), and PV(j)

It exposes more independent work, but FA3 Appendix B.3 reports two costs. First, relative to the two-stage design it needs another probability tile and row-scale state:

$$M_{\text{3-stage extra}} = B_rB_c\operatorname{sizeof}(\text{input}) + B_r\operatorname{sizeof}(\text{float}).$$

At B_r=B_c=128 with FP16 probabilities, this is 32.5 KiB, about 32.5 more logical 32-bit words per consumer thread for G=2, before layout overhead. Added to the 160-word two-stage subtotal, it leaves much less of a 240-register source budget for everything else.

Second, the paper reports that its compiler did not realize the intended three-way overlap: only the first WGMMA overlapped useful non-WGMMA work, while the second remained effectively unpaired. The paper therefore found the three-stage variant worse in practice. More graph-level parallelism is not useful when register pressure forces a smaller tile or lowering does not preserve the schedule.

Variant	Extra live-state obligation	Compiler obligation	Main failure mode
One stage	O plus current S/P transition	ordinary WGMMA correctness	scalar path serializes matrix work
Two stage	enforce S_next + P_current + O coexistence	QK/PV group age and scalar interleaving	spills or premature waits
Three stage	add next P and row scale beyond two stage	preserve three-way overlap	register pressure plus failed lowering

The source-reported FA3 Table 2 ablation also shows why the techniques cannot be treated as additive toggles. For non-causal FP16 forward at {batch,seqlen,nheads,hdim}={4,8448,16,128} on H100:

Paper configuration	Time	Throughput
full FA3	3.538 ms	661 TFLOP/s
warp specialization, no GEMM-softmax pipeline	4.021 ms	582 TFLOP/s
GEMM-softmax pipeline, no warp specialization	4.105 ms	570 TFLOP/s

These are source-reported measurements for one coordinate. The last two rows remove different components; their difference is not a clean estimate of either component’s isolated contribution.

The Compiler Schedule Needs A Proof Ladder

Source order is only the first rung. A release artifact should preserve this identity before claiming overlap:

source:
    FlashAttention release/commit and exact file hash

dependency closure:
    resolved CuTe DSL, CUTLASS, Quack, CUDA, nvcc, ptxas, target, flags

compile coordinate:
    dtype, B_r, B_c, d_h, G, stages, RS/SS PV path,
    rescale placement, causal/mask branch, generated kernel symbol

artifacts:
    IR/PTX hash, cubin hash, SASS hash, ptxas registers, local memory, SMEM

The proof then proceeds in four rungs:

1. Source DAG: QK commits before PV; wait_group(1) precedes score consumption; wait_group(0) precedes P reuse, V release, and post-PV O access.
2. PTX contract: wgmma.fence, two wgmma.commit_group operations, then wgmma.wait_group 1 and 0 appear in the legal regions. PTX can still differ in exact instruction count after specialization.
3. SASS static evidence: identify the corresponding WARPGROUP.ARRIVE, HGMMA, and dependency-barrier instructions; inspect instruction order, register count, and local-memory spill traffic. Static order supports a schedule argument but does not prove simultaneous engine activity.
4. Dynamic evidence: profile the exact symbol and coordinate, verify tensor/scalar activity, dependency and barrier stalls, achieved occupancy, local-memory traffic, correctness, and timing distribution.

FA3 Appendix B.2 supplies a source-reported example of rung 3 for the paper artifact. Its simplified SASS places row-max/exponential work early, interleaves F2FP, MUFU, and scalar instructions among the first HGMMA bundle, emits the second WGMMA bundle without comparable non-WGMMA interleaving, and later drains with a dependency barrier. This supports the paper’s two-stage explanation. It does not prove that the current beta23 source compiled in this repository produces the same SASS.

Compiler And Liveness Failure Matrix

Fault	Why source review may miss it	Observable evidence
helper turns wg_wait=-1 into a wait	dependency source still looks staged	PTX/SASS serializes after each GEMM
wait_group(0) emitted where 1 was intended	numerically correct	no scalar/PV overlap; higher dependency stalls
score read before the covering wait	high-level IR may obscure async use	invalid PTX contract or wrong results
P registers overwritten before PV completion	values appear logically “consumed” at issue	wrong output; register-input use before wait
O rescaled while PV is in flight	scalar statement looks independent	accumulator race/undefined behavior
K/V stage released before its WGMMA completes	pipeline counters remain balanced	producer overwrites data still being read
compiler extends S, P, or mask temporaries	equations count only payloads	high register count, lower residency, spills
SASS contains desired order but runtime does not overlap	static listing has no timing axis	profiler shows pipe or dependency stalls
profile captures another specialized kernel	symbol family looks similar	compile key/hash mismatch

The design rule is:

dependency graph exposes overlap
    -> wait placement makes it legal
    -> liveness determines whether it fits
    -> compiler artifacts show whether lowering preserved it
    -> profiler and correctness runs show whether it helped

0x06 FA3 FP8: Layout Is Part Of The Numerical Path

Changing Element from BF16 to FP8 is not a dtype substitution. It changes four contracts at once:

Contract	Choice made in the FA3 FP8 path	What breaks if it is hidden
representation	E4M3 operands and probability fragments, FP32 matrix accumulators	input width gets confused with accumulation or output width
layout	FP8 WGMMA shared-memory operands are k-major; P register ownership must match operand A	mathematically valid tiles reach the instruction in an illegal physical arrangement
scaling	Q/K factors enter scores; V factors enter output contributions; the axis determines legal factorization	a scale is applied after a nonlinear operation or broadcast over the wrong block
transform	the same orthogonal transform may precede Q/K quantization	pre-quantization equivalence is misreported as exact FP8 arithmetic

The owner question is therefore narrower than M8’s model-wide quantization question:

How does one FA3 Hopper forward path turn E4M3 Q/K/V and FP32 online-softmax state into legal FP8 WGMMAs without losing scale, reduction-axis, or layout semantics?

The answer needs two evidence coordinates. The FA3 paper specifies natural-block quantization and an incoherence transform. The pinned fa4-v4.0.0.beta23 C++ Hopper path at commit 2409214... exposes a different public scale granularity. They are related, not identical implementations.

E4M3 Names The Operand, Not The Whole Arithmetic Path

E4M3FN has one sign bit, four exponent bits, and three explicit fraction bits. Its largest finite magnitude is 448; its smallest positive subnormal is $2^{-9}=0.001953125$. The finite-only variant has a NaN encoding rather than infinity. A value outside the finite range must saturate, become non-finite under another conversion policy, or be prevented by scale selection.

The pinned public C++ admission path accepts PyTorch Float8_e4m3fn Q/K/V on SM90 or newer, requires the last dimension to be contiguous, requires head dimensions to be multiples of 16, and returns BF16 output for FP8 input. Internally, the Hopper template recognizes E4M3 and E5M2 as FP8 types, but that internal predicate does not widen this public API contract. Its forward path is:

Q, K, V storage:          E4M3FN
QK WGMMA accumulator S:  FP32
softmax row state:        FP32
probability operand P:   E4M3FN after FP32 exponentiation
PV WGMMA accumulator O:  FP32
public output:            BF16

This is mixed-precision attention. Calling it an “FP8 accumulator” would be wrong.

Derive Where Every Scale Must Enter

Let the real-valued blocks be approximated by E4M3 code values and positive descales:

$$Q_i \approx d_i^Q Q_i^{(8)}, \qquad K_j \approx d_j^K K_j^{(8)}, \qquad V_j \approx d_j^V V_j^{(8)}.$$

For query block $i$ and K/V block $j$, the unquantized score block is

$$S_{ij}=\alpha Q_iK_j^{\mathsf T}.$$

Substitution gives the low-precision matrix product and its dequantization factor:

$$\widehat S_{ij} = \alpha d_i^Qd_j^K \left(Q_i^{(8)}K_j^{(8)\mathsf T}\right).$$

The product $d_i^Qd_j^K$ must affect the score before row maximum, exponentiation, and normalization. Applying it after softmax is not equivalent because softmax is nonlinear.

The V scale has a different factorization. Partitioning the output over K/V traversal blocks gives

$$O_i = \sum_j P_{ij}V_j \approx \sum_j d_j^V \left(P_{ij}V_j^{(8)}\right).$$

If $d_j^V$ changes with j, each PV tile contribution must enter the online output accumulation in the right scale frame. One final scalar cannot be pulled outside the sum. If one $d^V_{b,h_{kv}}$ is constant for the whole batch/KV-head traversal, then

$$\widehat O_i = d^V_{b,h_{kv}} \sum_j P_{ij}V_j^{(8)},$$

and the factor can legally be folded into final row normalization.

This exposes a paper/source divergence that must remain visible:

Surface	Q descale axis	K descale axis	V descale axis	Consequence
FA3 Section 3.3	query block [B_r,d]	K block [B_c,d]	V block [B_c,d_v]	score and PV contribution factors can vary per traversal tile
pinned beta23 public C++ API	[batch, kv_head]	[batch, kv_head]	[batch, kv_head]	Q/K product is constant for that head; one V factor can be applied at finalize
missing descale pointer	implicit 1.0	implicit 1.0	implicit 1.0	caller has asserted that E4M3 values already use the intended numerical units

The pinned API checks exact descale shape (batch_size, num_heads_k), including Q descale. It does not accept the paper’s query-block or K/V-block scale arrays through this interface. The bundled FP8 benchmark also passes unit descales; that benchmark scaffolding is not evidence that natural-block scaling is implemented by the public call.

Why The Source Raises Softmax Probabilities By 2^8

Ordinary unnormalized softmax values lie in [0,1], using little of E4M3’s positive range and making small probabilities more likely to underflow. The pinned Hopper Softmax<..., Max_offset=8> instead forms

$$u_{ij}=2^8\exp(\widehat S_{ij}-m_i).$$

The FP32 row sum receives the same factor:

$$\ell'_i = \sum_j u_{ij} = 2^8\sum_j\exp(\widehat S_{ij}-m_i) = 2^8\ell_i.$$

Ignoring P conversion error for one moment,

$$\frac{\sum_j u_{ij}V_j}{\ell'_i} = \frac{2^8\sum_j\exp(\widehat S_{ij}-m_i)V_j}{2^8\ell_i},$$

so the offset cancels exactly. The source also divides the stored sum by 2^8 when reconstructing log-sum-exp. The numerical purpose is not to change attention; it is to present a wider [0,256] probability range to E4M3.

Conversion makes the actual path approximate:

$$\widehat O_i = \frac{d^V}{\ell'_i} \sum_j \mathcal Q_{\mathrm{E4M3}}(u_{ij}) V_j^{(8)}.$$

The denominator is accumulated from FP32 u, while PV consumes rounded E4M3 probabilities. The offset can reduce underflow; it does not remove E4M3 rounding or saturation.

Renderer-safe scale path:

raw QK WGMMA:
    C_ij = q8_i @ k8_j.T                  # FP32 accumulator

score entering softmax:
    S_ij = alpha * dQ_i * dK_j * C_ij

probability operand:
    u_ij = 2^8 * exp(S_ij - row_max_i)    # FP32
    p8_ij = E4M3(u_ij)                    # approximate

pinned head-wide V scale:
    O_i = dV_bh * sum_j(p8_ij @ v8_j) / sum_j(u_ij)

paper block-varying V scale:
    O_i = sum_j[dV_j * (p8_ij @ v8_j)] / sum_j(u_ij)

FP8 WGMMA Turns V Layout Into A Producer Job

For FP16/BF16 Hopper WGMMA, a shared-memory operand can use the relevant mn-major or k-major mode. FA3 notes that FP8 WGMMA accepts only k-major shared-memory operands. Q and K commonly arrive contiguous in head dimension, which matches their reduction dimension in QK. V also commonly arrives contiguous in head dimension, but in P @ V the reduction dimension is sequence/block position. TMA can tile and swizzle; it does not exchange the contiguous matrix dimensions.

Three designs are possible:

Design	Benefit	Cost or constraint
store V sequence-contiguous upstream	direct FP8 PV consumer layout	changes the external V/cache layout and every producer/consumer agreement
standalone global transpose	simple attention mainloop	another launch and HBM read/write unless fused elsewhere
in-kernel tile transpose	preserves head-contiguous public input and avoids standalone HBM traffic	extra SMEM buffer, producer instructions, barrier state, and shape constraints

FA3 and the pinned C++ path choose the third design for row-major V. The source sets Transpose_V = Is_FP8 && !V_colmajor, uses MN as V’s TMA major mode and K as its WGMMA major mode, and expands the producer from one warp to a full 128-thread warpgroup.

The source-level state machine is:

1. TMA writes head-contiguous V_j into Vt_smem[stage].
2. The producer warpgroup waits on pipeline_vt[stage].full.
3. LDSM.T loads packed 16-bit units from Vt_smem to registers.
4. Each 16-bit unit contains two FP8 values; byte_perm separates/reorders them.
5. STSM writes the k-major tile into V_smem[stage].
6. An async shared-memory fence publishes the transposed view.
7. pipeline_v[stage] is committed for the WGMMA consumer.
8. A 128-thread named barrier closes collective register/shared-memory use.
9. Only then is the original Vt stage released to TMA.

The audited templates require d_v % 32 == 0, B_c % 32 == 0, and either d_v % 64 == 0 or B_c % 64 == 0. They choose 64 x 32 or 32 x 64 transpose blocks accordingly. The separate Vt source and V destination stages are not free merely because FP8 uses one byte per value.

The Physical Transpose Includes A Paired Mathematical Permutation

LDSM.T and STSM are described around 16-bit entries. Packing two FP8 values into each entry leaves an 8-bit ordering problem that transpose instructions alone cannot solve. FA3 Appendix B.8 inserts byte permutations with selectors 0x6420 and 0x7531 between the collective load and store.

The faster pinned layout also permutes V’s eventual output-feature columns. The consumer must compensate. For a permutation $\sigma$ of the common reduction coordinate,

$$PV = P_{:,\sigma}V_{\sigma,:}.$$

This is an exact reindexing identity, not a quantization transform. Permuting only P or only V silently changes the result. The row-major-V path later calls permute_output_fp8 to undo its chosen output-column arrangement.

FP32 Score Ownership Is Not FP8 Operand-A Ownership

The first WGMMA writes S into FP32 accumulator registers. After exponentiation, those same logical entries must become the FP8 register-A operand of the second WGMMA. Downcasting values does not automatically move each byte to the lane and register expected by operand A.

FA3 Figures 3, 4, and 9 derive the four-lane exchange with __byte_perm and __shfl_sync. The pinned C++ source contains two specialized routes:

V route	P conversion/ownership repair	Output repair
head-contiguous V with in-kernel transpose	permute the FP32 C-register view, reinterpret it as A-register layout, then convert to E4M3	undo the chosen V-column permutation in O
already column-major V	reinterpret/convert, then use byte-permute and four-lane shuffle on E4M3 A registers	column-major-specific epilogue mapping

The hot path around the conversion is therefore tensor-level, not merely scalar:

// Pinned-source structure, abbreviated.
online_softmax(tSrS);                    // FP32 u, row_sum remains FP32
if (fp8 && row_major_v) permute_Cregs_fp8(tSrS);
tOrP_acc = reinterpret_as_pv_A_layout(tSrS);
convert_type_out(tOrP_acc, tOrP);        // FP32 -> E4M3
if (fp8 && v_colmajor) permute_Aregs_fp8(tOrP);
gemm_pv(tOrP, transposed_or_colmajor_V, tOrO); // FP32 O accumulator

Shape-compatible values in the wrong register ownership are still wrong. A test that checks only matrix dimensions cannot catch this class of bug.

This is a source-derived contract reconstruction from FA3 Sections 3.3, B.7, and B.8 plus the pinned beta23 Hopper C++ path. It is not a captured GPU trace.

Incoherent Processing Is Exact Before Quantization Only

Let $R$ be orthogonal:

$$RR^{\mathsf T}=I.$$

Apply the same right transform to Q and K:

$$Q'=QR, \qquad K'=KR.$$

Then

$$Q'K'^{\mathsf T} = QRR^{\mathsf T}K^{\mathsf T} = QK^{\mathsf T}.$$

FA3 follows the incoherence-processing lineage of QuIP and QuIP#, using products of random sign diagonals and a normalized Hadamard transform so multiplication costs $O(d\log d)$ rather than $O(d^2)$. The paper says the transform can be fused with an upstream memory-bound RoPE operation. Fusion means one kernel can perform both operations; it does not mean an arbitrary rotation commutes with RoPE.

It should not be summarized as the whole kernel “using QuIP.” FA3 adapts one incoherence-processing idea to its Q/K FP8 path; it does not import QuIP’s weight-quantization objective, codebooks, or complete method.

V is not multiplied by this Q/K orthogonal transform. Do not conflate it with the paired P/V layout permutation above. The former changes the head-feature basis before quantization while preserving QK in exact arithmetic. The latter only reindexes the PV reduction axis.

After E4M3 conversion,

$$\mathcal Q(QR)\mathcal Q(KR)^{\mathsf T} \ne QK^{\mathsf T}$$

in general. The hoped-for benefit is geometric: a signed Hadamard mixes a concentrated coordinate into many coordinates, often making one shared scale less destructive. It is not universal.

Executable Oracle: Exact Identities And Geometry-Dependent Error

The FP8 evidence bundle contains a pure-Python E4M3FN codebook oracle. It is seed-fixed and accumulates in Python binary64; it is not a bit-exact CUDA conversion or a local H100 run.

For seed 3406, it verifies:

Check	Maximum absolute error
$(QR)(KR)^{\mathsf T}=QK^{\mathsf T}$	1.78e-14
paired $P_{:,\sigma}V_{\sigma,:}=PV$	2.22e-16
exact 2^8 numerator/denominator cancellation before P rounding	1.11e-16

Its paper-inspired input distribution is $N(0,1)+N(0,100)\operatorname{Bernoulli}(0.001)$ over 32 x 32 Q/K/V fixtures. That run contains one injected outlier in each tensor:

CPU teaching variant	Output RMSE versus binary64 reference
per-tensor E4M3	0.04762
per-8-row-block E4M3	0.04066
signed Hadamard plus per-block E4M3	0.05024

The transform is worse on that seed. This is useful evidence: the identity is exact, but the quantized-error outcome is distribution- and scale-dependent.

A separate controlled fixture places persistent +/-10 excursions in Q/K channel 0. Per-block E4M3 RMSE is 0.23677; adding the same signed Hadamard to Q/K reduces it to 0.08066. That experiment illustrates the intended geometry without pretending to reproduce FA3 Table 3.

What The Paper Actually Reports

FA3 Section 4.3 uses an FP64 reference and the same stated synthetic distribution. Its Table 3 reports:

Paper method	RMSE
baseline FP16	3.2e-4
FA2 FP16	1.9e-4
FA3 FP16	1.9e-4
baseline FP8, per-tensor scaling	2.4e-2
FA3 FP8, block quantization + incoherent processing	9.1e-3
FA3 FP8 without block quantization	9.3e-3
FA3 FP8 without incoherent processing	2.4e-2

The often-quoted 2.6x is 2.4e-2 / 9.1e-3 for this source-reported coordinate. The ablation says that incoherent processing dominates this particular synthetic result and block quantization provides a smaller additional change here. It does not establish the same ordering for every model, layer, sequence, or training run. The paper also lists large-scale FP8 training behavior as future work rather than evidence already supplied.

Paper Algorithm, Pinned Source, And Runtime Selection

The complete claim ladder for this section is:

FA3 paper:
    natural-block scales, orthogonal Q/K transform,
    FP8 layout algorithms, source-reported accuracy and throughput

pinned beta23 C++ source:
    E4M3 public admission, [batch, KV-head] descales,
    Max_offset=8, V transpose, register permutations, BF16 output

compiled artifact:
    actual FP8 instruction variants, register allocation, SMEM,
    conversion behavior, selected layout specialization

runtime evidence:
    selected symbol, correctness against a reference,
    saturation/underflow counters, latency and throughput

Only the first two rungs are present here, plus a CPU oracle. The pinned artifact was not built, the FP8 kernel was not run locally, and no H100 profile was captured.

FP8 Failure Matrix

Fault	Why it can look plausible	Consequence
broadcast Q scale by KV block	block counts can accidentally match	wrong score scale for each query block
multiply $d^Qd^K$ after softmax	scale is still present somewhere in code	different distribution because softmax is nonlinear
apply one final V scale when $d_j^V$ varies	correct for a head-wide scale	cross-tile contributions use inconsistent units
feed head-contiguous V directly to FP8 PV WGMMA	logical V shape is correct	illegal or semantically wrong major-mode layout
release Vt before collective transpose closes	destination stage may look ready in one warp	producer overwrite or partial transposed tile
permute P without matching V, or V without matching P	each tensor remains in bounds	exact PV result changes silently
convert FP32 P without register ownership repair	element count and dtype match	wrong lanes/bytes feed WGMMA operand A
add 2^8 to P but not its denominator/LSE bookkeeping	larger P range appears intentional	output or log-sum-exp is scaled incorrectly
zero, negative, NaN, or infinite descale	pointer shape passes API checks	collapse, NaN propagation, or overflow
let one scale drive values beyond 448	E4M3 conversion still returns a byte	saturation destroys magnitude information
say the orthogonal transform makes FP8 exact	exact identity was proven	quantization errors are omitted from the claim
report one favorable transform seed only	deterministic result is reproducible	geometry dependence is hidden by selection

FP8 Boundary And Release Gate

This section owns the FA3-specific low-precision kernel path: its E4M3 operands, FP32 state, scale application points, WGMMA layout repair, V transpose, and exact identities surrounding quantization.

M8.3 owns general outlier geometry, equivalent-transform topology, conditioning, and method comparison. M8.6 owns model-wide live activation and K/V scale lifetimes, append state, cache metadata, and serving integration. M4.6 does not turn this one attention path into a general quantization policy.

M8.2 owns model-wide calibration, sensitivity, and precision allocation; M8.9 owns quantized artifact compatibility and runtime admission; M8.10 owns deployment evaluation, promotion, release, and rollback of the exact candidate. The checklist below concerns this article’s local FP8 explanation; it does not authorize deployment of a quantized candidate.

Publish only if:

• E4M3 input/P width is separated from FP32 accumulation and BF16 output.
• every descale names its axis and application point.
• paper block scales are separated from pinned [batch, kv_head] API scales.
• the 2^8 probability offset is derived through numerator, denominator, and LSE.
• V transpose includes source/destination stages, packed-FP8 byte movement, synchronization, and shape constraints.
• accumulator-to-operand register ownership is traced after conversion.
• Q/K orthogonal equivalence and paired P/V permutation are proved separately.
• exact pre-quantization identities are not promoted to exact E4M3 arithmetic.
• paper Table 3 remains source-reported and separate from the CPU oracle.
• no local GPU accuracy, throughput, or compiled-layout claim is implied.

0x07 Blackwell Changes The Machine Again

The FlashAttention-4 paper targets B200/GB200 Blackwell. Its algorithm is not FA3 with larger constants. Blackwell changes where matrix accumulators live, who may issue the matrix operation, how a CTA pair shares work, and which synchronization event makes a result consumable. A schedule that preserves the algebra but violates one of those machine contracts is not a slower FA4 kernel. It is an invalid kernel.

This loop asks a deliberately narrower owner question:

What must be true of TMEM addresses, lifetimes, CTA-group ownership, and asynchronous completion before an FA4 schedule can be legal?

The next section derives the forward schedule and its rate model. This section first fixes the machine on which that schedule must execute.

Resource	Hopper / FA3 model	Blackwell / FA4 model	Kernel consequence
Tensor MMA	WGMMA	tcgen05.mma / UMMA	new instruction and synchronization contract
Accumulator location	registers	TMEM	accumulator lifetime moves out of general registers
Cooperative MMA	one warpgroup-oriented path	1-CTA and 2-CTA modes	paired CTAs can share a larger MMA tile
Representative accumulator tile	64 x 128 in FA3 comparison	128 x 128	row ownership and softmax mapping change
BF16 MMA throughput in FA4 model	4096 ops/cycle/SM	8192 ops/cycle/SM	matrix work doubles per cycle
MUFU exponential rate in FA4 model	16 ops/cycle/SM	16 ops/cycle/SM	softmax becomes relatively more expensive
SMEM read rate in FA4 model	128 B/cycle/SM	128 B/cycle/SM	shared-memory feeding can become limiting

The rate comparison is the FA4 paper’s simplified per-SM model. It identifies design pressure; it is not a complete GPU roofline and excludes resources such as register bandwidth, L2, instruction issue, and occupancy. The paper’s main text and performance figures say B200, while Appendix A.1 says B100 and contains a date inconsistent with the arXiv submission. This article preserves that metadata discrepancy instead of silently choosing one label.

TMEM Adds A New State Coordinate

The PTX ISA describes Tensor Memory, or TMEM, as a dedicated on-chip store distinct from registers and shared memory. tcgen05.mma can place its accumulator there. A later tensor operation can consume a TMEM fragment directly, while ordinary elementwise work first loads the required TMEM lanes into registers.

That gives one logical value at least four independent coordinates:

logical:       score S, probability P, or output accumulator O
address:       TMEM lane and column
time:          allocation epoch and producer/consumer lifetime
ownership:     CTA group, CTA rank, warp, and lane slice

S0 and P0 may refer to the same columns at different times. Two CTAs may use the same numeric TMEM address in different per-CTA views. One warp may issue an instruction whose result is not yet visible to another warp. A scalar pointer alone therefore does not identify the state.

Derive Capacity, Granularity, And Address

For sm_100, PTX exposes a per-CTA addressable view with 512 columns, 128 lanes per column, and one 32-bit cell at each lane-column coordinate. The FA4 paper describes the corresponding 256 KiB capacity in its per-SM machine model:

$$M_{\mathrm{TMEM}} =512\ \text{columns} \times128\ \text{lanes} \times4\ \text{bytes}$$ $$M_{\mathrm{TMEM}} =262{,}144\ \text{bytes} =256\ \text{KiB}.$$

The minimum allocation granule is 32 columns:

$$M_{32} =32\times128\times4 =16{,}384\ \text{bytes} =16\ \text{KiB}.$$

The legal allocation amounts are powers of two between 32 and 512 columns:

32, 64, 128, 256, 512 columns

Successive allocation amounts in one kernel may not increase, and all allocated TMEM must be explicitly deallocated before kernel exit. Calling relinquish_alloc_permit closes the right to allocate again; it does not make live values disappear.

A 32-bit TMEM address packs lane and column:

$$\operatorname{taddr}(\lambda,c) = (\lambda\ll16)\;|\;c, \qquad 0\le\lambda<128, \quad 0\le c<512.$$

For example, lane 37, column 384 encodes as 2,425,216, then decodes back to (37, 384). This is an address inside one CTA’s view. Multiplying it by the CTA count would invent a flat address space that PTX does not provide.

The warp index determines which lane band a TMEM load or store owns:

Warp inside warpgroup	TMEM lanes
0	[0, 32)
1	[32, 64)
2	[64, 96)
3	[96, 128)

One warp covers 32 lanes. A warpgroup of four warps is needed to cover all 128 lanes. This is different from the issue rule of tcgen05.mma, which is a single-thread operation.

Allocation Is A Collective Lifetime

The pinned fa4-v4.0.0.beta23 Blackwell forward path allocates the maximum 512 columns. Its control path is structurally:

# Pinned-source structure, abbreviated.
if warp_idx == mma_warp_id:
    tmem.allocate(get_max_tmem_alloc_cols("sm_100"))  # 512 columns
    tmem.wait_for_alloc()
    tmem_ptr = tmem.retrieve_ptr(Float32)
    run_mma_pipeline(...)
    tmem.relinquish_alloc_permit()
    tmem_alloc_barrier.arrive_and_wait()
    tmem.free(tmem_ptr)

The allocation wrapper hides different hardware participation contracts:

Operation family	CTA group 1 issue	CTA group 2 issue	Completion/visibility obligation
allocate, deallocate, relinquish	one warp in the CTA	one warp from each peer CTA	collective operation must complete before dependent use or exit
mma, cp, shift, commit	one thread	one thread from the CTA pair	result is asynchronous
ld, st, wait	warp operation	warp operation in each relevant CTA	lane ownership and operation-specific wait apply

“Single-thread issue” does not mean “single-thread computation.” The one issuing thread launches a collective tensor operation whose operands and destination are distributed according to the MMA atom and CTA-group contract.

For CTA group 2, the peer CTA must already be active and remain active while a paired operation can touch its state. The peer rank differs in the final cluster-rank bit. Allocation, relinquish, and deallocation require one participating warp from each CTA. Deallocating after only the local CTA has stopped touching TMEM is insufficient; peer access must also have quiesced.

Derive The Pinned Beta23 TMEM Map

The current beta23 source does not leave the offsets implicit. For a K/V tile width $N$, value head dimension $d_v$, two score slots, and $q_{\mathrm{stage}}\in\{1,2\}$ output slots, it constructs:

$$c_{S_0}=0, \qquad c_{S_1}=N,$$ $$c_{O_i}=2N+i d_v, \qquad 0\le i<q_{\mathrm{stage}},$$

so the logical map requires

$$C_{\mathrm{logical}} =2N+q_{\mathrm{stage}}d_v \le512.$$

With $N=128$, $d_v=128$, and two Q stages:

Object	Columns	Width	Live use
FP32 score S0	[0, 128)	128	QK accumulator, masking, softmax input
BF16 probability P0	[64, 128)	64	PV operand after S0 dies
FP32 score S1	[128, 256)	128	second score/softmax slot
BF16 probability P1	[192, 256)	64	PV operand after S1 dies
FP32 output O0	[256, 384)	128	persistent PV accumulator
FP32 output O1	[384, 512)	128	second persistent PV accumulator

The source represents offsets in FP32-column units. Since BF16 P has half the width of FP32 S, one probability slot needs

$$C_P = \left\lceil N\frac{16}{32} \right\rceil =\frac N2 =64$$

columns for this coordinate. The P slot is placed in the tail half of the dead S slot. This is not “P and S coexist in 64 columns.” It is a time-dependent alias.

The capacity equation also exposes why stage count and output width cannot be tuned independently:

N=128, q_stage=2, d_v=128 -> 2*128 + 2*128 = 512 columns: exact fit
N=128, q_stage=1, d_v=192 -> 2*128 + 1*192 = 448 columns: algebraically fits
N=128, q_stage=2, d_v=192 -> 2*128 + 2*192 = 640 columns: impossible

The last two rows are sensitivity probes, not claims that the pinned dispatcher selects both coordinates. The source requests all 512 columns even when its logical map uses fewer. Its fragment code can treat base address zero as known precisely because a full-width allocation starts at zero; a smaller allocation would need a retrieved-base-dependent layout.

This is an editorial reconstruction from PTX ISA 9.3, the official CUTLASS tcgen05 guide, and beta23 commit 2409214.... It is not a captured B200 trace.

Alias Legality Requires Address And Time

For objects $x$ and $y$, define their column intervals $C_x,C_y$ and live intervals $L_x,L_y$. Reuse is legal only when at least one coordinate does not overlap:

$$\operatorname{legal\_alias}(x,y) \iff (C_x\cap C_y=\varnothing) \lor (L_x\cap L_y=\varnothing).$$

The score-to-probability transition is therefore a protocol:

tcgen05.mma writes FP32 S_i to TMEM
-> completion is observed
-> softmax warps load their owned S_i lanes
-> all score consumers finish
-> BF16 P_i overwrites the tail of the same address interval
-> PV consumes P_i as a TMEM operand
-> P_i dies before the score slot is reused

Changing P0.time_begin from the score-release boundary to an earlier instant produces a real live-address conflict even though every address remains in bounds. Capacity checks alone cannot detect this bug.

Async Issue, Completion, And Visibility Are Different Events

tcgen05.mma is asynchronous. The issuing thread can continue before the TMEM result is safe for a consumer load. The positive dependency chain is:

single thread issues tcgen05.mma
-> single thread commits completion to an mbarrier
-> consumer observes the mbarrier completion
-> participating threads synchronize as required
-> tcgen05.fence::after_thread_sync publishes the TMEM result
-> four warps load their 32-lane slices
-> tcgen05.wait::ld closes the load before reuse

The exact PTX spellings depend on instruction form, but the state transitions do not. A barrier arrival is not itself a consumer fence, and a consumer fence cannot precede completion observation.

Shared-memory operands add a second memory-proxy boundary. Generic-proxy writes and asynchronous-proxy tensor operations that touch the same SMEM location require the appropriate fence.proxy.async ordering. TMEM completion does not retroactively publish an unfenced SMEM producer.

PTX permits pipelined same-thread MMA ordering only under its matching-accumulator, shape, and kind conditions. Changing the MMA kind or accumulator relationship removes that implicit dependency; the schedule must then add explicit ordering. “Issued by the same thread” is not a universal happens-before proof.

One CTA And Two CTAs Partition Different Coordinates

The official CUTLASS tcgen05 programming guide distinguishes CTA-group 1 and CTA-group 2 atoms. In the pinned source:

cta_group_size = 2 if use_2cta_instrs else 1
mma_tiler_qk = (cta_group_size * M, N, d)
mma_tiler_pv = (cta_group_size * M, d_v, N)
cluster_shape_mn = (cta_group_size, 1)

For $M=N=d=128$:

Contract	CTA-group 1	CTA-group 2
logical MMA tile	128 x 128 x 128	256 x 128 x 128 across the pair
CTA 0 output rows	[0, 128)	[0, 128)
CTA 1 output rows	separate tile	[128, 256) in the paired tile
cluster shape	(1, 1)	(2, 1)

The pair expands the M coordinate. It does not automatically split the reduction K coordinate, and it does not merge the two per-CTA TMEM views into one flat 1024-column allocation.

The B operand is logically needed by both output-row owners. Two independent CTAs would each stage a full BF16 128 x 128 B tile:

$$M_B=128\times128\times2=32\ \text{KiB},$$ $$M_{B,\,\mathrm{independent\ pair}}=2M_B=64\ \text{KiB}.$$

In the pinned 2-CTA layout, the source divides K/V SMEM per stage by the CTA-group size and uses cluster-aware TMA/UMMA pipelines. Each CTA stages a 16 KiB share, so the pair stages 32 KiB in this teaching coordinate:

$$\frac{M_{B,\,\mathrm{2CTA\ pair}}} {M_{B,\,\mathrm{independent\ pair}}} = \frac{32}{64} =0.5.$$

This is a source-derived shared-memory payload ledger, not a universal claim that 2-CTA mode halves HBM traffic or doubles end-to-end speed. TMA multicast, DSMEM access, pair synchronization, cluster residency, tails, and occupancy still decide the realized cost.

A CTA Pair Is Not Two Independent CTAs

The paired mode adds contracts that do not exist between ordinary independent blocks:

1. Both CTAs must select CTA-group 2; all tcgen05 operations in one kernel must use the same CTA-group mode.
2. The peer CTA must be active before a paired operation and may not exit while the pair can still access its TMEM state.
3. Each CTA selects its own MMA slice with get_slice(mma_tile_coord_v) and writes its own output portion.
4. Pipeline participant counts that bridge UMMA and non-MMA work scale by cta_group_size.
5. Allocation, relinquish, and deallocation are pair-collective even though mma issue is single-thread.
6. Cluster placement is a feasibility constraint. It must be checked with the architecture-appropriate cluster occupancy API, not inferred from thread count alone.

The Blackwell tuning guide reports 64K 32-bit registers per SM, a maximum of 255 registers per thread, 228 KiB of shared memory per SM, and 227 KiB per CTA. Those ceilings bound admission; they do not prove that a 16-warp specialization, its SMEM, and a two-CTA cluster achieve any particular occupancy. Only a compiled artifact and cluster occupancy query can close that claim.

Pinned Dispatch Makes Two-CTA Conditional

The current beta23 public interface does not select the general 2-CTA path for every Blackwell attention call. At commit 2409214..., the ordinary path requires all of the following:

SM100/SM110 family
2CTA not explicitly disabled
noncausal and nonlocal attention
no split-KV output
fixed-length Q (no cu_seqlens_q or seqused_q)
no block sparsity
page size absent or 128
padded d in {128, 192}
padded d_v == 128
packed query length > 2 * tile_m
GQA packing divisibility condition satisfied

A separate dedicated path admits (d,d_v)=(256,256). These predicates are pinned implementation behavior, not requirements imposed by the FA4 paper or the tcgen05 ISA. A later release may change them. The only runtime proof that 2-CTA executed is the selected compiled symbol/configuration plus a profile; source eligibility alone is not dispatch evidence.

Executable TMEM And CTA-Group Oracle

The Blackwell TMEM evidence bundle turns the contracts above into a deterministic CPU model. It regenerates the 512-column source map, address round trip, legal allocation sizes, live-alias audit, 1-CTA/2-CTA ownership, B-payload ledger, and positive completion trace.

Its negative fixtures fail with named reasons:

Mutation	Required failure
allocate 96 columns	allocation is not a legal power of two
allocate 64 then 128 columns	allocation size increases
allocate after relinquishing	allocation permit is closed
use two Q stages with algebraic $d_v=192$	logical map needs 640 columns
write P while aliased S is still live	simultaneous address/lifetime overlap
mix CTA-group 1 and 2 operations	kernel-wide CTA-group mode mismatch
launch group 2 without active peer	peer-liveness violation
allocate group 2 from one warp only	collective-participant violation
load TMEM before completion/fence	unfinished result consumption
deallocate before peer quiescence	live peer access
let warp 0 load lane 64	lane-ownership violation

This oracle checks arithmetic and event preconditions. It does not lower PTX, prove ptxas instruction selection, run on B200, or measure cluster occupancy.

Blackwell Failure Matrix

Fault	Why it can look plausible	Consequence
treat TMEM as extra registers	both hold accumulator-like values	wrong allocation and visibility protocol
count a CTA pair as one 1024-column pool	each CTA exposes 512 columns	invalid cross-CTA address arithmetic
overlap P with a still-live S tail	both intervals fit inside 512 columns	softmax input is overwritten before all consumers finish
request 512 columns, then allocate again after relinquish	first allocation remains live	illegal allocation-permit transition
let one CTA allocate/deallocate group 2	MMA issue itself is single-thread	collective allocation protocol is incomplete
allow peer CTA to exit early	local CTA still has valid pointers	paired instruction accesses dead peer state
load after mbarrier wait but before consumer fence	completion was observed	TMEM visibility to the consumer is not established
use warp 0 for all 128 lanes	pointer arithmetic is in range	96 lanes are outside that warp’s ownership
mix CTA-group modes in one kernel	each instruction form is individually valid	kernel violates a global tcgen05 constraint
assume same-thread MMA issue always orders operations	source order looks sequential	unsupported accumulator/shape/kind transition races
call 2-CTA universally faster	source halves one SMEM payload ledger	cluster, synchronization, tail, or occupancy cost can dominate
treat source eligibility as runtime selection	predicate evaluates true on paper	compiled specialization or dispatcher may differ

Blackwell Hardware Boundary And Release Gate

This section owns the Blackwell machine contract needed before FA4 scheduling: TMEM shape and addressing, allocation lifetime, source-derived suballocation, time-safe aliasing, tcgen05 issue/completion rules, CTA-group ownership, and pinned 2-CTA dispatch admission.

It does not yet claim that the FA4 forward schedule hides exponential latency or reaches the paper’s throughput. Section 0x08 owns the forward pipeline and feeds-and-speeds derivation. Section 0x09 owns software exponential and conditional rescaling. Section 0x0A owns backward, DSMEM reduction, and determinism.

Publish this loop only if:

• 256 KiB capacity is derived without inventing one flat group address space.
• allocation size, permit, deallocation, lane, and CTA-group rules are separate.
• the beta23 S/P/O map is derived from $N$, $q_{\mathrm{stage}}$, and $d_v$.
• address overlap is legal only when live intervals do not overlap.
• single-thread issue is separated from asynchronous completion and visibility.
• group-2 output ownership is separated from K-reduction ownership.
• the 0.5 B-payload ratio is labeled source-derived, not a universal speedup.
• current 2-CTA dispatch predicates are pinned behavior, not paper or ISA requirements.
• the CPU oracle is not presented as a compiled Blackwell artifact or local benchmark.
• FA4 forward scheduling, exponential implementation, and backward DSMEM remain deferred.

This changes the capacity problem:

Hopper:
    many live accumulator tiles consume registers directly

Blackwell:
    accumulator tiles consume finite TMEM columns
    registers still hold softmax rows and scalar/vector state
    SMEM still stages matrix operands

Moving an accumulator does not eliminate its lifetime. It gives the scheduler a different storage pool and a different access path.

2-CTA MMA

Blackwell can pair two CTAs for a cooperative MMA. Each CTA stages part of the operands and owns an accumulator slice, while the instruction operates as one larger tile. The CTAs form a fixed cooperating pair and must agree on the paired mode and instruction order. This can reduce duplicated shared-memory reads, but it adds cluster placement, distributed shared-memory exchange, and pairwise synchronization constraints.

0x08 FA4 Forward: Deriving The Feeds And Speeds

The owner question for this loop is narrower than “is attention compute-bound?”:

For one Blackwell forward work tile, what work feeds the tensor, shared-memory, and exponential paths; which dependencies permit that work to overlap; and which source-level schedule attempts to approach the resulting lower bound?

The answer requires three different objects. The FA4 paper supplies a deliberately simplified machine model and an intended two-query-tile schedule. The official machine contracts in 0x07 constrain what TMEM and tcgen05 are allowed to do. The pinned fa4-v4.0.0.beta23 source supplies concrete warp IDs, pipeline objects, split-$P$ handoffs, and startup/steady/drain order. None of those is a measured runtime trace.

This section owns:
    forward work and resource-time derivations
    instruction-ceiling sensitivity
    paper H/L query-tile schedule
    beta23 warp roles and pipeline handoffs
    startup, steady-state, and drain ordering
    q_stage and TMEM-partition alternatives

It consumes:
    0x07 TMEM allocation, alias, and CTA-group legality
    M4.3 online-softmax algebra
    M4.4 work-partition language

It defers:
    software exp2 accuracy and conditional rescaling -> 0x09
    backward, DSMEM, and deterministic reduction -> 0x0A
    measured throughput and dispatch proof -> 0x0D and profiler evidence

Asymmetric Scaling Changes What Deserves A Ledger

The paper’s B200/GB200 model uses three per-SM rates:

Resource	Paper-model rate	Evidence and limit
BF16 tensor operations	$R_{\mathrm{MMA}}=8192$ FLOPs/cycle/SM	derived in the paper from its cited B200 peak, clock, and SM count
SMEM reads	$R_{\mathrm{SMEM}}=128$ bytes/cycle/SM	source-reported microbenchmark rate, not an ISA guarantee
exponential operations	$R_{\exp}=16$ operations/cycle/SM	paper-reported B200/GB200 MUFU rate

The paper explicitly calls this a simplified analysis. It omits ordinary FP32/FMA work, register and TMEM movement, barrier latency, issue limits, L2/HBM effects, occupancy, tails, and epilogue cost. Its purpose is to expose pressure, not predict a kernel timestamp.

That limitation is central to the history. If tensor throughput doubles while SMEM and exponential throughput do not, the old statement “the tensor cores dominate” can become false without the attention equation changing. A pipeline redesign is then an accounting response to asymmetric hardware scaling.

Derive The Two MMA Terms From Tensor Shapes

Let one logical forward tile cover:

Q: [M, d]
K: [N, d]
V: [N, d]
S = Q K^T: [M, N]
P = softmax(S): [M, N]
O = P V: [M, d]

For a conventional multiply-add count, $[M,d]\times[d,N]$ performs $MNd$ multiplies and $MNd$ additions:

$$F_{QK^\top}=2MNd.$$

The second contraction has $M\times N\times d$ multiply-adds as well:

$$F_{PV}=2MNd.$$

Therefore the tensor-work ledger is

$$F_{\mathrm{MMA}}=F_{QK^\top}+F_{PV}=4MNd,$$

and the paper-model resource time is

$$T_{\mathrm{MMA}} = \frac{4MNd}{R_{\mathrm{MMA}}}.$$

This is logical work divided by a peak per-SM rate. It does not yet include partial-tile efficiency, instruction latency, or inability to keep the tensor path continuously issued.

Derive SMEM Reads From MMA Instruction Coverage

The SMEM term is not “read Q, K, and V once.” The paper models 128-by-128 MMA output tiles and counts repeated operand reads when $M$, $N$, or $d$ requires multiple instructions.

For $QK^\top$, both operands come from SMEM. Define

$$I_{QK} = \left\lceil\frac{M}{128}\right\rceil \left\lceil\frac{N}{128}\right\rceil.$$

Each covered output tile reads one $128\times d$ Q operand and one $d\times128$ K operand. With $b$ bytes per element,

$$B_{QK} = b I_{QK}(128d+128d).$$

For $PV$, $P$ is a TMEM operand and V is the modeled SMEM operand. Its instruction grid is

$$I_{PV} = \left\lceil\frac{M}{128}\right\rceil \left\lceil\frac{d}{128}\right\rceil,$$

so the V-read ledger is

$$B_{PV} = b I_{PV}(128N).$$

The general paper-model time is therefore

$$T_{\mathrm{SMEM}} = \frac{B_{QK}+B_{PV}}{R_{\mathrm{SMEM}}}.$$

When $M$, $N$, and $d$ are divisible by 128, the ceilings disappear. Substitution gives

$$B_{QK}=\frac{bMNd}{64}, \qquad B_{PV}=\frac{bMNd}{128},$$

and hence

$$T_{\mathrm{SMEM}} = \frac{3bMNd}{128R_{\mathrm{SMEM}}}.$$

For BF16, $b=2$, and for the paper’s $128$ bytes/cycle/SM rate,

$$T_{\mathrm{SMEM}} = \frac{3MNd}{8192}.$$

This is SMEM-to-tensor-operand traffic after global tiles have been staged. It is not the HBM traffic owned by M4.3, and it is not the full on-chip traffic of the compiled kernel.

Derive Exponential Pressure And The Crossover Head Dimension

There is one exponentiation per score entry in the simplified model:

$$N_{\exp}=MN, \qquad T_{\exp}=\frac{MN}{R_{\exp}}.$$

The ratio to tensor work removes $M$ and $N$:

$$\frac{T_{\exp}}{T_{\mathrm{MMA}}} = \frac{R_{\mathrm{MMA}}}{4dR_{\exp}}.$$

The two resource times are equal at

$$d^* = \frac{R_{\mathrm{MMA}}}{4R_{\exp}}.$$

With $R_{\mathrm{MMA}}=8192$ and $R_{\exp}=16$, $d^*=128$. This is why the canonical head-dimension-128 ledger ties MMA and exponentiation. It is not a universal property of attention: change the head dimension or either hardware rate and the balance moves.

For divisible BF16 tiles, the SMEM/MMA ratio is also dimensionless:

$$\frac{T_{\mathrm{SMEM}}}{T_{\mathrm{MMA}}} = \frac{3bR_{\mathrm{MMA}}}{512R_{\mathrm{SMEM}}} = 0.75.$$

The last equality is specific to the paper’s rates and BF16 element width.

max Is A Lower Bound, Not A Runtime Formula

If the three modeled resources could be overlapped perfectly, a lower bound would be

$$T_{\mathrm{modeled,lb}} = \max(T_{\mathrm{MMA}},T_{\mathrm{SMEM}},T_{\exp}).$$

If none overlapped, the modeled subtotal would be

$$T_{\mathrm{modeled,serial}} = T_{\mathrm{MMA}}+T_{\mathrm{SMEM}}+T_{\exp}.$$

Neither expression is the complete kernel runtime. The first ignores dependency bubbles and all omitted resources. The second includes no omitted resources either and therefore is not a rigorous upper bound. The useful question is how close a legal schedule can move the modeled subtotal toward the modeled lower bound.

For $(M,N,d)=(128,128,128)$:

$$T_{\mathrm{MMA}}=1024, \qquad T_{\mathrm{SMEM}}=768, \qquad T_{\exp}=1024.$$

The lower bound is 1024 cycles and the no-overlap subtotal is 2816 cycles. Calling the kernel a “1024-cycle kernel” would be wrong; 1024 is only the largest term in this simplified resource ledger.

Repository-authored analytic reconstruction from FA4 Section 3.1.1 and Table 1. The diagram is not a benchmark or profiler capture.

Instruction Ceilings Explain Ragged-Tile Surprises

The compact $3MNd/8192$ SMEM formula is legal only under the paper’s divisibility assumption. Consider $(M,N,d)=(192,96,160)$. Logical MMA and exponential work scale with the requested shape, but the modeled instruction grids are

QK instruction grid: ceil(192/128) x ceil(96/128) = 2 x 1
PV instruction grid: ceil(192/128) x ceil(160/128) = 2 x 2

The oracle obtains:

Tile $(M,N,d)$	MMA cycles	SMEM cycles	exp cycles	Largest modeled term
$(128,128,128)$	1024	768	1024	MMA and exp tie
$(256,128,128)$	2048	1536	2048	MMA and exp tie
$(192,96,160)$	1440	2048	1152	SMEM after instruction ceilings

The ragged row is a teaching coordinate, not a claim that beta23 selects that exact tile. It shows why dropping ceilings can name the wrong resource before code generation begins.

Sensitivity Moves The Bottleneck Without Changing The Algorithm

Holding the canonical 128-cubed shape fixed:

Synthetic rate mutation	MMA	SMEM	exp	Largest term
paper B200 model	1024	768	1024	MMA and exp
double only exponential throughput	1024	768	512	MMA
double only tensor throughput	512	768	1024	exp
halve only SMEM bandwidth	1024	1536	1024	SMEM

This table is a sensitivity analysis, not a forecast for another GPU. It demonstrates the design principle: optimizing the currently largest term can expose the next one, and a tile choice should be re-evaluated against the actual target’s rates and instruction shapes.

The Paper Schedule Uses Two Query Tiles To Create Independent Work

The FA4 paper’s Figure 1 names two 128-row query/output tiles H and L. The intended pipeline performs tensor work for one tile while the other tile’s softmax warpgroup processes a score row. Blackwell changes the ownership geometry relative to FA3:

Hopper FA3:
    accumulator in registers
    multiple threads participate in one logical score row

Blackwell FA4 paper schedule:
    accumulator in TMEM
    one 128-thread softmax warpgroup
    each thread owns one complete 128-value score row

One-thread-per-row removes the inter-thread row-max shuffle required by the Hopper ownership pattern, but it does not make softmax cheap. A thread must hold the 128 score values and temporary state in registers. The paper cites 128 registers for the input row and potentially 64 for output-related state, before miscellaneous temporaries. Register redistribution is therefore part of the schedule’s feasibility.

The two softmax warpgroups are not two independent exponential engines. The paper explicitly synchronizes them so their exponential critical sections do not overlap. Their other work can overlap tensor operations, but the scarce MUFU interval remains serialized.

Source: Zadouri et al., FlashAttention-4, Figure 1, CC BY-NC-SA 4.0. Superscripts H and L identify the two 128-row query/output tiles. This is the paper schedule, not a complete beta23 source trace or measured timeline.

Two S Slots Win The Startup Trade-Off

After reserving two output tiles in TMEM at head dimension 128, the paper describes two broad partitions for the remaining space:

TMEM alternative	Benefit	Cost
one S slot + two independent P slots	simpler non-aliased P placement	only one score tile can occupy the remaining score surface at startup
two S slots whose tails become P	compute both startup score tiles immediately	requires address-plus-lifetime proof for the S/P aliases

FA4 chooses the second. Section 0x07 derived the corresponding beta23 map: S0/P0, S1/P1, O0, and O1. The purpose of the overlap is now visible: two complete score tiles can be issued before the main PV/QK alternation begins.

Beta23 Turns The Paper Roles Into Sixteen Warp IDs

The pinned flash_fwd_sm100.py partitions one CTA as follows for the audited two-stage path:

Warp IDs	Source role	State emphasis
0..3	softmax stage H	one 128-row score tile, large row-register budget
4..7	softmax stage L	second score tile, same logical role
8..11	correction	consume scale statistics, conditionally rescale O, final normalization/epilogue path
12	MMA issuer	allocate TMEM; issue QK and PV operations; own tensor order
13	epilogue	separate S-to-G path when correction warps do not own it
14	load	Q/K/V producer path when one load warp is sufficient
15	empty or CLC scheduler	specialization-dependent spare/scheduler role

These IDs are pinned implementation behavior, not requirements of the FA4 algorithm. The source can reassign roles: q_stage=1 removes the second softmax warpgroup, non-TMA K/V paths can consume two load warps, and the correction warps can absorb epilogue work.

For one audited noncausal BF16, $d=128$, 1-CTA teaching specialization, the source fallback targets are

softmax setmaxnreg target:   192
correction setmaxnreg target: 80
other-role target:            48

budget equation:
    2 * 192 + 80 + 48 = 512

Those are setmaxnreg policy targets in source. They are not compiler-reported registers per thread, spill counts, or an occupancy result. A compiled cubin and toolchain-specific resource report are still required.

Pipeline Objects Encode The Actual Handoffs

The paper’s arrows become several distinct beta23 pipelines around lines 928-1026:

Source object	Producer -> consumer	Payload/lifetime meaning
pipeline_q	load -> MMA	Q stage is full; MMA later releases it after the K/V traversal
pipeline_kv	load -> MMA	one circular K or V stage is full; release permits producer reuse
pipeline_s_p_o	MMA <-> softmax + correction	MMA publishes S; softmax publishes P and correction publishes rescaled O before PV/reuse
pipeline_p_lastsplit	softmax -> MMA	final P fraction becomes available after the early split signal
pipeline_o_acc	MMA -> correction	final accumulated O is available when a separate guarantee is required
pipeline_sm_stats plus named barrier	softmax -> correction	row-max/row-sum scale state is ready
pipeline_o_epi	correction -> epilogue	corrected output is ready for a separate output-store role

pipeline_s_p_o is intentionally not a normal one-way queue. The MMA warp first acts as producer of S. The softmax and correction roles then close the stage from the opposite side by producing P readiness and O-rescale readiness. Replacing that circuit with one vague “barrier” would hide which state is safe to reuse.

Startup, Steady State, And Drain Are Different Schedules

For two query stages and two K/V blocks, the pinned MMA loop’s logical tensor-issue order is

startup:
    QK_H_0
    QK_L_0

steady alternation:
    PV_H_0 -> QK_H_1
    PV_L_0 -> QK_L_1

drain:
    PV_H_1
    PV_L_1

The source startup waits for Q and the first K stage, issues both score operations, publishes both S slots, and only then enters the K/V loop. The steady loop waits for P plus corrected O, issues PV, then overwrites the now-dead S/P slot with the next QK result. The drain releases Q and issues the final PV operations without another QK behind them.

This is why a steady-state diagram is not a full algorithm. A short K/V traversal spends a larger fraction of time in startup and drain, and empty/sparse work must still close the synchronization protocol without fabricating payloads.

Repository-authored reconstruction from FA4 Figure 1 and pinned beta23 source lines 263-337, 928-1026, 1193-1303, 1555-1847, and 2254-2553. It is not a captured Blackwell timeline.

Split P Creates An Issue Point, Not An Early Completion

Beta23 sets

$$C_{P,\mathrm{early}} = 32\left\lfloor\frac{3N/4}{32}\right\rfloor.$$

For $N=128$, this is 96 columns, or 75 percent of P. The source comment and constructor explain the intent: softmax stores the first fraction, signals the MMA warp so PV can begin, then stores and signals the remainder.

The dependency is therefore two-level:

P_75_ready + corrected_O_ready + V_ready
    -> PV may issue

P_100_ready
    -> PV may consume the final P fraction and complete

PV complete
    -> aliased S/P columns may become the next S destination

Treating the 75-percent signal as “P is complete” would permit an early overwrite or an incomplete contraction. The split is overlap within one logical PV operation, not permission to weaken its completion contract.

q_stage, Paper $M$, And CTA-Group $M$ Are Distinct Coordinates

The paper’s Table 1 uses $M=256$ for two 128-row query tiles. In the pinned source, three different coordinates contribute to covered query rows:

$$M_{\mathrm{work\ group}} = q_{\mathrm{stage}} \times g_{\mathrm{CTA}} \times M_{\mathrm{block}}.$$

With $M_{\mathrm{block}}=128$:

Source configuration	Query rows in one logical work group	TMEM columns for $N=d_v=128$	Softmax warpgroups
q_stage=1, g_CTA=1	128	$2N+d_v=384$	1
q_stage=2, g_CTA=1	256	$2N+2d_v=512$	2
q_stage=2, g_CTA=2	512 across the CTA pair	two per-CTA 512-column views	2 per CTA

The interface pin selects q_stage=2 when the packed query length exceeds one M tile and otherwise uses one stage, with dedicated paths adding further constraints. Paper $M=256$ must not be copied into every 2-CTA source ledger as though it named the same coordinate.

Design Alternatives Preserve Algebra But Change The Schedule

Choice	What improves	What worsens or changes
one query stage	less TMEM and one softmax warpgroup	less independent H/L work; startup cannot precompute two S slots
two query stages	more overlap opportunity and paper-like H/L pipeline	fills 512 TMEM columns at $N=d_v=128$; two large register owners
one S + two P regions	simpler P address lifetime	weaker two-S startup
two S regions aliased with P	immediate dual-QK startup	requires exact P/PV completion before reuse
unsplit P publication	simpler readiness protocol	PV waits for all P columns
75/25 split publication	overlaps final P production with PV issue	adds partial/full completion states
rescale O in softmax role	fewer role handoffs	lengthens softmax critical work and register lifetime
separate correction warpgroup	removes O rescale from the softmax critical path	adds statistics, O-ready, and epilogue coordination

These are schedule alternatives, not alternative attention equations. A valid implementation can choose differently if its resource rates, head dimensions, or feature path differ.

Source-Level Forward Skeleton

# Teaching skeleton. Pipeline objects and phases are omitted, not optional.
load_Q_stages()
load_first_K_stage()

for q_stage in (H, L):
    wait(Q[q_stage], K[0])
    issue_QK(Q[q_stage], K[0], S[q_stage])
    publish_S_full(q_stage)

for kv_block in remaining_blocks:
    for q_stage in (H, L):
        # These are separate arrivals in the real source circuit.
        wait(P_early[q_stage])
        wait(O_rescaled[q_stage])
        wait(V[kv_block - 1])
        issue_PV(P[q_stage], V[kv_block - 1], O[q_stage])
        wait(P_final[q_stage])
        complete_PV(q_stage)

        # S and P share columns; overwrite only after PV consumption closes.
        wait(K[kv_block])
        issue_QK(Q[q_stage], K[kv_block], S[q_stage])
        publish_S_full(q_stage)

for q_stage in (H, L):
    wait(P_early[q_stage], P_final[q_stage], O_rescaled[q_stage], V[last])
    issue_and_complete_final_PV(q_stage)
    publish_final_O(q_stage)

The actual source pipelines K and V through the same circular SMEM surface, carries phase state across persistent work tiles, handles masks and sparse/empty paths, and can fuse correction with the epilogue. The skeleton exists to expose dependencies, not replace the implementation.

Executable Forward Oracle

The forward-schedule evidence bundle contains:

loop8_forward_schedule_fixture.json
loop8_forward_schedule_oracle.py
loop8_forward_schedule_report.json
loop8_forward_schedule_source_trace.json

The CPU oracle mechanically:

1. reproduces both FA4 Table-1 cycle rows;
2. retains MMA instruction ceilings for a ragged teaching tile;
3. derives the $d^*=128$ MMA/exp crossover;
4. verifies that the 16 source roles partition warp IDs 0..15;
5. checks the 96/128 split-P fraction and source register-target equation;
6. validates startup/steady/drain tensor order and required handoff edges;
7. rejects overlapping normalized H/L exponential critical sections;
8. rejects nine resource, dependency, alias, and evidence-surface mutations.

It does not compile CuTe DSL, lower tcgen05, simulate B200 pipelines, establish occupancy, or measure overlap.

Forward Failure Matrix

Tempting shortcut	Why it appears plausible	Broken contract
drop instruction ceilings	logical shapes are smaller than covered tiles	repeated full-operand reads disappear from the SMEM ledger
add resource times and call that optimized runtime	all three terms are cycles	overlap has been forbidden by assumption
take max and call it measured runtime	perfect overlap lower bound is compact	dependencies and omitted resources disappear
overlap H and L exponential critical sections	there are two softmax warpgroups	the paper deliberately serializes scarce MUFU work
issue PV before correction releases O	P is already available	old O may still be in the wrong exponential frame
overwrite S after the early P signal	75 percent of P is published	PV still needs final P and the alias remains live
equate paper $M=256$ with every source work group	both discuss larger query coverage	q_stage and CTA-group M were collapsed
call setmaxnreg(192) the compiled register count	source contains the number	compiler allocation and spills were never inspected
call the source issue order a profiler timeline	order is explicit in Python DSL	lowering, issue gaps, stalls, and runtime selection are unknown

Renderer-Safe Formula Ledger

F_MMA = 4 * M * N * d
T_MMA = F_MMA / R_MMA

I_QK = ceil(M/128) * ceil(N/128)
B_QK = b * I_QK * (128*d + 128*d)

I_PV = ceil(M/128) * ceil(d/128)
B_PV = b * I_PV * (128*N)
T_SMEM = (B_QK + B_PV) / R_SMEM

T_exp = M * N / R_exp
d_crossover = R_MMA / (4 * R_exp)

T_modeled_lower_bound = max(T_MMA, T_SMEM, T_exp)

Forward-Schedule Boundary And Release Gate

Publish this loop only if:

• the two MMA terms are derived from shapes rather than asserted;
• SMEM reads retain 128-by-128 instruction ceilings before simplification;
• exponential pressure is tied to $MN$ and a hardware-rate coordinate;
• max(resource times) is labeled an optimistic lower bound, not runtime;
• paper H/L roles are separated from beta23 warp IDs;
• H/L exponential critical sections do not overlap in the paper contract;
• split-P issue, final-P completion, and S/P alias reuse are separate events;
• startup, steady state, and drain are all represented;
• paper $M$, q_stage, and CTA-group M are not collapsed;
• register targets are not presented as compiler artifacts;
• 0x09 still owns exponential approximation and conditional-rescale numerics;
• no local Blackwell build, profile, or timing is implied.

0x09 FA4 Softmax: Use More Than The Exponential Unit

The forward ledger in 0x08 ended with a simplified B200 model in which a 128 x 128 x 128 forward tile needs 1024 modeled tensor cycles and 1024 modeled exponential cycles. Faster tensor cores therefore do not make the scalar path disappear. They make the scalar path visible.

The owner question for this loop is:

How can one softmax warpgroup feed more exponentials than one saturated special-function path can provide, while preserving a coherent online-softmax reference and a finite-precision probability range?

The answer has two coupled parts. The FA4 paper moves a subset of 2^x evaluations onto FP32 FMA/ALU resources. It also delays some output-accumulator rescalings so the correction warpgroup performs fewer vector passes over O. The pinned fa4-v4.0.0.beta23 source makes both choices compile-time specialization parameters rather than universal constants.

This section owns:
    base-2 score, reference, and probability coordinates
    degree-3 exp2 range reduction and Horner evaluation
    hardware/software exp2 dispatch within one score row
    arbitrary-reference numerator/denominator recurrence
    threshold, offset, warp-vote, and correction-warpgroup contracts
    deterministic CPU error and recurrence oracles

It consumes:
    M4.3 exact online-softmax and merge algebra
    0x06 probability-offset and FP8 range coordinates
    0x07 TMEM/correction ownership
    0x08 H/L warp roles, split-P, and resource pressure

It defers:
    backward dS exponential and reduction schedule -> 0x0A
    measured specialization speedups -> 0x0D and profiler evidence
    model-wide calibration, sensitivity, and precision allocation -> M8.2
    live activation quantization and scale/state lifecycle -> M8.6

One Score Has Raw, Natural, And Base-2 Coordinates

Let $s$ denote the score value held after QK, masking, and score modification, before the ordinary softmax scale $\sigma$. The mathematical softmax exponent is

$$z = \sigma s.$$

The source evaluates exp2, so define

$$\lambda = \sigma\log_2 e, \qquad b = \lambda s.$$

If $a$ is the retained score reference in the same coordinate as $s$, and $\omega$ is an optional probability offset, the actual exp2 argument is

$$x = \lambda(s-a)+\omega = b-r+\omega, \qquad r=\lambda a.$$

These coordinates must not be collapsed:

Coordinate	Meaning	Pinned source representation
$s$	score before softmax scale	FP32 score fragment loaded from TMEM
$a$	retained score reference	SoftmaxSm100.row_max[0]
$b=\lambda s$	base-2 logit	produced by packed FMA in scale_subtract_rowmax
$r=\lambda a$	base-2 accumulator reference	implicit in row_max * scale_log2
$x=b-r+\omega$	exponent argument	consumed by hardware or emulated exp2
$2^x$	unnormalized probability operand	converted to BF16/FP16/FP8 before PV

This is also where the conditional-rescale threshold lives. A source threshold $\tau=8$ means

$$(a_{\text{candidate}}-a_{\text{retained}})\lambda \le 8,$$

not “the raw logits differ by at most eight” unless $\lambda=1$.

Derive The Software exp2 Path

For one base-2 exponent argument $x$, Cody-Waite-style range reduction writes

$$n=\lfloor x\rfloor, \qquad f=x-n\in[0,1),$$

and therefore

$$2^x=2^n2^f.$$

The pinned helper approximates only the fractional term. Its degree-3 coefficients, stored from lowest to highest power, are:

$$\begin{aligned} p_0 &= 1,\\ p_1 &= 0.6951461434364319,\\ p_2 &= 0.22756439447402954,\\ p_3 &= 0.07711908966302872. \end{aligned}$$

Thus

$$2^f \approx p_0+p_1f+p_2f^2+p_3f^3,$$

but the implementation does not materialize powers. Horner’s rule gives

$$p(f)=p_0+f\left(p_1+f\left(p_2+fp_3\right)\right),$$

which needs three dependent FP32 FMAs per scalar. evaluate_polynomial_2 evaluates two score elements with packed f32x2 FMAs.

The integer path is less obvious. The source:

1. clamps $x$ from below at -127;
2. adds $C=2^{23}+2^{22}=12,582,912$ with round-toward-minus-infinity;
3. subtracts $C$ using the expected round-to-nearest mode, recovering $n=\lfloor x\rfloor$;
4. computes $f=x-n$;
5. evaluates $p(f)$;
6. shifts the encoded integer bits into the FP32 exponent position and combines them with the bits of $p(f)$.

The magic constant works because the FP32 unit in that magnitude range is one. The directed-rounding addition removes the fractional part while leaving the small integer encoded in the low bits used by the subsequent exponent-field construction. The source assumes $x\le127$; the softmax reference and dtype envelope keep the positive side far below that bound.

# Teaching equivalent of utils.ex2_emulation_2; source uses packed FP32x2 ops.
def exp2_degree3(x):
    x = max(x, -127.0)
    n = floor_via_directed_rounding_magic(x)
    f = x - n

    fractional = fma(p3, f, p2)
    fractional = fma(fractional, f, p1)
    fractional = fma(fractional, f, p0)

    return combine_exponent_bits(n, fractional)

This is not a general replacement for a correctly rounded exponential library. The lower clamp deliberately makes every $x<-127$ share the same effective input; the bit construction assumes a bounded positive side; and the coefficient/error contract is specialized to exp2 on the reduced interval.

Repository-authored reconstruction from FA4 Section 3.1.3 and pinned beta23 softmax.py/utils.py. It is an algorithm and source map, not captured SASS or a throughput trace.

Partial Emulation Is A Fragment Predicate, Not One Percentage

The paper reports emulating roughly 10-25% of entries in its tested configurations. The pinned source exposes the actual selection rule in SoftmaxSm100.apply_exp2_convert.

For a row fragment of 32 elements, the loop visits element pairs $(k,k+1)$. A pair uses software emulation only when all three conditions hold:

$$k\bmod q \ge q-r_e,$$ $$j\ge j_{\text{start}}, \qquad j<F-1,$$

where $q$ is ex2_emu_freq, $r_e$ is ex2_emu_res, $j$ is the fragment index, and $F$ is the number of 32-element fragments. The last fragment always uses hardware exp2; early fragments may also be forced to hardware.

For a 128-element teaching row:

Pinned specialization coordinate	freq/res/start	Emulated entries	Fraction
SM100 2-CTA, noncausal, $d=128$	10/4/1	24 / 128	18.75%
SM100 1-CTA, causal, $d=128$	16/4/1	16 / 128	12.50%
SM100 2-CTA, noncausal, $d=192$	16/4/0	24 / 128	18.75%
SM100 1-CTA, causal, $d=192$	32/4/1	8 / 128	6.25%
SM100 2-CTA, $d=256$	14/6/0	36 / 128	28.125%
SM103 tuning rows	0/4/0	0 / 128	0%

The last two rows are especially important. They show later pinned-source tuning outside the paper’s compact 10-25% description. They do not contradict the paper experiment; they belong to a different release/configuration coordinate. ex2_emu_freq=0 means all hardware exp2, and the source comment attributes that choice to faster native exp2 on SM103.

Full software emulation is not automatically faster. It adds coefficient and intermediate live ranges, packed FMA traffic, integer/bit operations, and a dependency chain longer than one hardware MUFU.EX2. The useful design point balances two queues without making registers or ordinary issue bandwidth the next bottleneck.

Raw FP32 Error And BF16-Consumed Error Are Different Claims

FA4 Table 2 reports four million random inputs on $[0,1)$ against an FP64 reference:

Path	FP32 max rel. error	FP32 mean rel. error	BF16 max rel. error	BF16 mean rel. error
ideal FP64 rounded to BF16	-	-	3.89e-3	1.41e-3
hardware MUFU.EX2	1.41e-7	3.04e-8	3.89e-3	1.41e-3
degree 3	8.77e-5	5.43e-5	3.90e-3	1.41e-3
degree 4	3.05e-6	1.84e-6	3.89e-3	1.41e-3
degree 5	1.44e-7	5.48e-8	3.89e-3	1.41e-3

The softmax-numerics CPU oracle evaluates the pinned degree-3 coefficients with FP32 FMA rounding on a separate deterministic 65,536-point uniform grid. It obtains:

raw FP32 max relative error:       8.7683755e-5
raw FP32 mean relative error:      5.4271432e-5
post-BF16 max relative error:      3.8975323e-3
post-BF16 mean relative error:     1.4095395e-3
exact BF16-code match:             99.0051%
within one BF16 code step:         100%

The close agreement is a useful independent check of the coefficients and evaluation order. It is not a reproduction of the paper’s random sample, GPU instruction, or SASS path. The result supports the paper’s narrow argument: when probabilities are consumed at BF16 precision, BF16 rounding dominates this degree-3 approximation on the tested domain. It does not make the raw FP32 approximation equivalent to hardware, nor does it establish end-to-end model accuracy.

Conditional Rescaling Is An Arbitrary-Reference Recurrence

M4.3 derived online softmax with the running maximum as its reference. This section needs a more general state because a skipped update deliberately keeps an older reference.

Work in base-2 coordinates. After some blocks, let:

$$r = \text{retained base-2 reference},$$ $$D = \sum_i 2^{b_i-r+\omega},$$ $$N = \sum_i 2^{b_i-r+\omega}v_i.$$

$D$ is the denominator represented in frame $r$; $N$ is the unnormalized output numerator in the same frame. The attention output is $N/D$.

For the next score block $B$, define the candidate

$$c=\max\left(r,\max_{b\in B} b\right).$$

The threshold policy commits a new reference only when

$$c-r>\tau.$$

Thus

$$r'= \begin{cases} c, & c-r>\tau,\\ r, & c-r\le\tau, \end{cases} \qquad \gamma=2^{r-r'}.$$

Both branches are one recurrence:

$$D'=\gamma D+\sum_{b\in B}2^{b-r'+\omega},$$ $$N'=\gamma N+\sum_{b\in B}2^{b-r'+\omega}v.$$

If the reference is retained, $r'=r$ and $\gamma=1$: old state is untouched and the new block is centered on the retained reference. If the reference changes, the same $\gamma$ converts both old numerator and denominator into the new frame.

The proof is direct:

$$\gamma D = 2^{r-r'}\sum_i2^{b_i-r+\omega} = \sum_i2^{b_i-r'+\omega},$$

and the same identity holds termwise for $N$. Therefore

$$\frac{N'}{D'} = \frac{\sum_i2^{b_i-r'+\omega}v_i} {\sum_i2^{b_i-r'+\omega}},$$

so the retained reference and the offset cancel in exact arithmetic.

This corrects a tempting but inaccurate description: after a skipped update, the source’s row_max is not necessarily the literal true running maximum. It is a legal retained frame. After the first three blocks in the deterministic fixture, the true base-2 maximum is 2, while the retained reference is still 0; normalization still matches direct attention.

Final Normalization And LSE Use The Retained Frame

Because $D$ includes the probability offset $\omega$, the natural-log sum-exp is

$$\operatorname{LSE} = \left(r+\log_2D-\omega\right)\ln2.$$

In the default source coordinate $r=a\lambda$, this becomes

$$\operatorname{LSE} = \left(a\lambda+\log_2D-\omega\right)\ln2,$$

which is the form used by the pinned epilogue. No separate “true maximum” is required: the denominator is interpreted in the same retained frame that produced it.

The final output scale is simply

$$D^{-1},$$

with the V descale applied where required. Again, numerator and denominator must refer to the same frame. Rescaling only one breaks the ratio.

Source Trace: Softmax Warpgroup To Correction Warpgroup

The pinned forward path encodes the recurrence across two specialized owners:

Event	Softmax owner	Correction/MMA consequence
load/reduce S	softmax warpgroup loads one row and computes candidate maximum	no O mutation yet
threshold decision	update_row_max* returns retained row_max and acc_scale	acc_scale=1 means no frame conversion
publish scale	softmax row writes acc_scale to sScale and signals stats barrier	correction warpgroup can inspect the row
center and exponentiate	packed FMA forms $b-r+\omega$; hardware/software predicate computes P	P is converted and stored to TMEM
split-P release	first configured columns can release early issue permission	does not complete P or close the alias
denominator update	row_sum <- acc_scale * row_sum + sum(P)	denominator remains in the retained frame
warp decision	correction warp executes vote_ballot_sync(scale < 1)	any row requiring correction activates the warp path
O conversion	rows multiply O by their own scale; scale-one rows remain unchanged	MMA receives “O rescaled” permission
final epilogue	reciprocal row sum normalizes O; optional LSE uses retained frame	output and statistics close consistently

The ballot is not merely a performance detail. Per-lane control flow around a warp-cooperative TMEM copy/rescale path would be unsafe. The source uses one warp-level decision; rows that skipped the reference update carry scale one, so participating in the common path leaves them unchanged.

# Source-level teaching skeleton; barriers and TMEM layouts remain mandatory.
candidate = max(retained_reference, rowmax(score_fragment))
delta_log2 = (candidate - retained_reference) * softmax_scale_log2

if delta_log2 > rescale_threshold:
    next_reference = candidate
    old_state_scale = exp2(-delta_log2)
else:
    next_reference = retained_reference
    old_state_scale = 1.0

x = score_fragment * softmax_scale_log2 \
    - next_reference * softmax_scale_log2 \
    + max_offset
probability = mixed_hardware_and_degree3_exp2(x)

publish_scale_to_correction(old_state_scale)
row_sum = old_state_scale * row_sum + sum(probability)

# In the dedicated correction warpgroup:
if warp_any(old_state_scale < 1.0):
    output_accumulator *= old_state_scale
publish_output_rescaled()

Repository-authored algebra/source trace from FA4 Section 3.1.4 and pinned beta23 SoftmaxSm100, forward softmax, and correction loops. It is not a profiler timeline.

Threshold And Probability Offset Form One Range Contract

The source does not use one threshold for every dtype. It also adds a probability offset for FP8 so values are represented higher in the finite FP8 range before PV. If the retained reference is stale by at most $\tau$, the largest unnormalized probability is bounded by

$$P_{\max}\le2^{\tau+\omega}.$$

Pinned dtype coordinate	$\tau$	$\omega$	Derived $P_{\max}$	Format maximum used by source comment
BF16/FP16 input path	8	0	256	at least FP16 65504
FP8 E4M3FN	4	4	256	448
other 8-bit path / E5M2 coordinate	4	8	4096	57344

For E4M3FN, the headroom ratio is only 448/256=1.75. This explains why threshold and offset cannot be tuned independently. Raising either can saturate converted P even though the numerator/denominator algebra remains exact over real numbers. Section 0x06 owns the broader FP8 layout and scaling argument; this section owns the local online-softmax envelope.

Finite precision also creates a lower-side risk: entries far below the reference underflow in P. Retaining an old reference does not make that risk worse. The retained $r$ sits at most $\tau$ below the running maximum, so each new exponent $b-r+\omega$ is up to $\tau$ larger than in a frame that always follows the maximum. A very negative entry must therefore lie further below the maximum before it underflows; the price is upper headroom, which the bound $P_{\max}\le2^{\tau+\omega}$ above already accounts for. Moving the reference too frequently adds more accumulator multiplications and rounding. The threshold balances range, work, and error; it is not a universal softmax constant.

Design Alternatives And Their Costs

Choice	Benefit	Cost / proof obligation
all hardware MUFU.EX2	shortest implementation, strong raw FP32 approximation	can saturate the exponential path exposed in 0x08
all degree-3 software emulation	removes MUFU demand	maximizes FMA/ALU, register, and dependency-chain pressure
partial fragment striping	overlaps MUFU and FMA resources	tuning is shape/architecture-specific; fragment exclusions matter
degree 4 or 5	lower raw FP32 error	one or two extra dependent FMAs and longer live ranges
always update reference	strongest centering and simplest invariant	more O rescale passes and correction synchronization
delayed update	fewer O conversions	larger P dynamic range and threshold/offset proof
correction in softmax warpgroup	fewer role handoffs	lengthens the softmax critical path and register liveness
dedicated correction warpgroup	decouples O rescale from exponential work	adds TMEM ownership and barrier protocol
per-row divergent correction	avoids scale-one multiplies in an abstract scalar loop	incompatible with warp-cooperative source path; use warp vote

Executable Softmax-Numerics Oracle

The evidence package includes:

loop9_softmax_numerics_fixture.json
    paper accuracy coordinates
    pinned coefficients, source hashes, and tuning predicates
    dtype threshold/offset envelopes
    exact-reference recurrence fixture
    negative mutation registry

loop9_softmax_numerics_oracle.py
    FP32 FMA Horner simulation
    BF16 round-to-nearest-even comparison
    source fragment-dispatch counter
    arbitrary-reference numerator/denominator recurrence
    LSE and dtype-envelope checks

loop9_softmax_numerics_report.json
    deterministic generated result

The oracle proves arithmetic and source-coordinate consistency. It does not reproduce MUFU.EX2, compile CuTe DSL, report register allocation, execute a Blackwell kernel, or measure a speedup.

Softmax-Numerics Failure Matrix

Symptom	Likely violated contract	First check
degree-3 path looks as accurate as hardware in raw FP32	BF16-consumed and raw approximation claims were collapsed	compare before conversion
emulation fraction differs from a paper percentage	source fragment/frequency specialization changed	inspect freq/res/start and forced hardware fragments
output drifts only when reference update is skipped	new block and old accumulator use different references	log retained reference, candidate, and exponent bias
denominator is plausible but output is wrong	only D or only N was frame-converted	compare identical $\gamma$ on both states
LSE is biased by a constant	max_offset omitted from final log correction	check $r+\log_2D-\omega$
FP8 path saturates P	threshold-plus-offset exceeds dtype envelope	compute $2^{\tau+\omega}$
correction path deadlocks/diverges	per-lane branch replaced the warp vote	inspect ballot and barrier participation
SM103 slows after copied SM100 tuning	emulation treated as architecture-independent	verify native-exp2 specialization
article claims measured gain	CPU oracle/source inspection crossed its evidence ceiling	require compiled kernel and profile

Renderer-Safe Formula Ledger

base-2 score:
    lambda = softmax_scale * log2(e)
    b = lambda * score

exp2 range reduction:
    n = floor(x)
    f = x - n in [0, 1)
    2^x = 2^n * 2^f

degree-3 Horner polynomial:
    p(f) = p0 + f * (p1 + f * (p2 + f * p3))

conditional retained reference:
    candidate = max(retained_reference, block_max)
    reference' = candidate if candidate-reference > tau else reference
    gamma = 2^(reference-reference')

same-frame recurrence:
    D' = gamma * D + sum(2^(b-reference'+offset))
    N' = gamma * N + sum(2^(b-reference'+offset) * value)

final normalization and LSE:
    output = N' / D'
    LSE = (reference' + log2(D') - offset) * ln(2)

probability range:
    P_max <= 2^(tau + offset)

Softmax-Numerics Boundary And Release Gate

Publish this loop only if:

• raw scores, natural logits, base-2 logits, retained references, and offsets are distinct;
• the degree-3 polynomial is tied to pinned coefficients and FP32 FMA order;
• range reduction, directed rounding, clamp, and exponent reconstruction are explained;
• paper random-sample errors and the deterministic CPU grid remain separate evidence;
• raw FP32 and post-BF16 errors are not conflated;
• partial emulation is derived from the beta23 fragment predicate;
• source tuning drift beyond 10-25% is explicit;
• delayed rescaling updates numerator and denominator in one retained frame;
• retained row_max is not mislabeled as the true running maximum after a skip;
• final output and LSE are derived from the retained reference and offset;
• threshold and FP8 offset share one finite-range envelope;
• correction uses a warp vote and a dedicated ownership handoff;
• 0x0A still owns backward softmax-gradient scheduling;
• no CPU oracle is presented as compiled Blackwell accuracy or speed.

0x0A FA4 Backward: Five MMAs, TMEM Lifetimes, And Ordered Reductions

Owner question. How does the attention-backward graph become a legal Blackwell kernel when five matrix contractions, softmax-gradient work, CTA-pair exchange, and global gradient accumulation all compete for different storage and synchronization domains?

This section owns the FA4 backward dataflow and its 1CTA and 2CTA realization. It consumes the exact softmax algebra from M4.3, the FA2 backward partition from M4.4, and the Blackwell machine contract from 0x07. It does not reopen FA1 IO complexity, teach tcgen05 from first principles, or own persistent work scheduling and benchmark reconciliation; those remain in 0x0B-0x0D.

The evidence has four levels and they are not interchangeable:

Evidence	What it establishes	What it does not establish
FA4 Sections 3.2.1-3.2.4	backward graph, analytic cycle ledger, 2CTA exchange, deterministic-order intent	the exact beta23 specialization or a measured B200 timeline
official NVIDIA documentation	CTA-group MMA, TMEM ownership, DSMEM lifetime, remote-copy byte accounting	which FlashAttention path was selected
pinned fa4-v4.0.0.beta23 source	concrete warp roles, TMEM offsets, pipeline objects, relay barriers, reduction calls	generated SASS, achieved occupancy, or progress for every specialization
deterministic CPU oracle	algebra, repacking, interval, event-order, and failure invariants	compiled Blackwell correctness or performance

Deriving The Backward Graph Instead Of Memorizing It

For one attention head, let

$$S=\alpha QK^\top,\qquad P=\operatorname{softmax}_{\text{row}}(S),\qquad O=PV,$$

where $Q\in\mathbb{R}^{M\times d}$, $K,V\in\mathbb{R}^{N\times d}$ for the paper’s equal-value-dimension model, and $\alpha$ is the score scale. Given $dO$, the linear map $O=PV$ gives two gradients immediately:

$$dV=P^\top dO,\qquad dP=dO V^\top.$$

The softmax step is the part that is easy to compress too far. For one row $p=\operatorname{softmax}(s)$,

$$\frac{\partial p_j}{\partial s_k}=p_j(\mathbf{1}_{j=k}-p_k).$$

Multiplying that Jacobian by the upstream row $dp$ gives

$$ds_j =\sum_k dp_k\frac{\partial p_k}{\partial s_j} =p_j\left(dp_j-\sum_k p_kdp_k\right).$$

Define the row scalar

$$\delta_i=\sum_j P_{ij}dP_{ij}.$$

Then the matrix form is

$$dS=P\odot(dP-\delta[:,\mathrm{None}]).$$

Because $O_i=\sum_j P_{ij}V_j$ and $dP_{ij}=dO_i\cdot V_j$, the same scalar can be obtained without materializing all of $P\odot dP$:

$$\delta_i =\sum_jP_{ij}(dO_i\cdot V_j) =dO_i\cdot\left(\sum_jP_{ij}V_j\right) =dO_i\cdot O_i.$$

Finally, differentiating $S=\alpha QK^\top$ yields

$$dQ=\alpha dS K,\qquad dK=\alpha dS^\top Q.$$

This derivation gives two useful invariants. First, every row of $dS$ sums to zero:

$$\sum_jdS_{ij} =\sum_jP_{ij}dP_{ij}-\delta_i\sum_jP_{ij} =0.$$

Second, delta = row_sum(P * dP) and delta = row_dot(dO, O) must agree. The backward oracle checks both identities before it tests any physical schedule.

Why The Paper Counts Five MMAs

Backward does not retain the full score or probability matrix from forward. It reconstructs them tile by tile. The matrix work is therefore:

1. S  = alpha * Q @ K^T       # recompute scores
2. dP = dO @ V^T              # upstream probability gradient
3. dV += P^T @ dO             # value gradient partial
4. dK += alpha * dS^T @ Q     # key gradient partial
5. dQ += alpha * dS @ K       # query gradient partial

P  = exp(S - LSE)             # elementwise reconstruction
dS = P * (dP - delta)         # elementwise softmax backward

The five MMAs are S, dP, dV, dK, and dQ. Reconstructing P and forming dS are elementwise stages between MMA producers and consumers. A source implementation may transpose or permute these tensors to satisfy tensor-operation layouts; those physical orientations do not change the logical equations.

Object	Logical tile shape	Produced by	Consumed by	Precision boundary
S	M x N	Q @ K^T MMA	exponential / P reconstruction	FP32 accumulator in TMEM
P	M x N	elementwise exp(S-LSE)	dV MMA and dS	source-specific operand conversion
dP	M x N	dO @ V^T MMA	dS	FP32 accumulator before elementwise work
dS	M x N	P * (dP-delta)	dK, dQ MMAs	converted/staged as an MMA operand
dV	N x d	P^T @ dO MMA	global reduction/epilogue	FP32 accumulator
dK	N x d	dS^T @ Q MMA	global reduction/epilogue	FP32 accumulator
dQ	M x d	dS @ K MMA	staged global reduce-add	FP32 accumulator

Source: Zadouri et al., FlashAttention-4, Figure 2. The figure is the paper schedule, not a beta23 instruction trace.

Reconstructing The Paper’s One-CTA Roofline

One tile performs five contractions, each with $2MNd$ floating-point operations under the usual multiply-add-equals-two convention:

$$F_{\mathrm{MMA}}=5(2MNd)=10MNd.$$

At the paper model’s normalized rate of 8192 operations per cycle,

$$T_{\mathrm{MMA}}=\frac{10MNd}{8192}.$$

The shared-memory operand ledger can also be derived rather than copied. Three MMAs use shared/shared-style operand traffic in the simplified model:

S:   Q (M*d)  + K  (N*d)
dP: dO (M*d)  + V  (N*d)
dQ: dS (M*N)  + K  (N*d)

The dV and dK contractions consume their TMEM-side accumulator operand and each read one M*d shared-memory operand. The total operand elements are therefore

$$(2Md+3Nd+MN)+2Md=4Md+3Nd+MN,$$

and the paper’s 64-element-per-cycle model gives

$$T_{\mathrm{SMEM,MMA}}=\frac{4Md+3Nd+MN}{64}.$$

Two additional transfers matter. dS is written as an MMA operand,

$$T_{dS,\mathrm{write}}=\frac{MN}{64},$$

and FP32 dQ is written to and read from shared memory before the global reduction,

$$T_{dQ,\mathrm{write+read}}=\frac{Md}{16}.$$

The elementwise exponential work is

$$T_{\exp}=\frac{MN}{16}.$$

For M=N=d=128, the complete normalized ledger is:

Resource component	Derivation	Cycles
five MMAs	10*M*N*d / 8192	2560
shared MMA operands	(4*M*d + 3*N*d + M*N) / 64	2048
dS write	M*N / 64	256
dQ FP32 write plus read	M*d / 16	1024
total shared-memory ledger	2048 + 256 + 1024	3328
exponential unit	M*N / 16	1024

The model takes the maximum resource time, not the sum:

T_tile_lower_bound ~= max(T_MMA, T_SMEM_total, T_exp)
                   = max(2560, 3328, 1024)
                   = 3328 cycles

Shared memory exceeds the MMA budget by

$$\frac{3328-2560}{2560}=30\%.$$

This is a peak-rate lower bound. It is not a profiler timeline and does not include launch, dependency bubbles, barriers, tails, atomics, or occupancy loss. Its value is causal: it explains why reducing shared-memory movement can matter even though backward contains more MMA work than forward.

From The Paper Graph To Beta23 Warp Ownership

The pinned flash_bwd_sm100.py specializes the graph into 16 warps, or 512 threads, per CTA:

Warp IDs	Source role	Principal ownership
0-3	reduction	dQ TMEM to registers to SMEM to global reduce-add
4-11	compute	reconstruct P, form dS, stage peer exchange, dK/dV epilogue
12	MMA	issue the five tcgen05 MMA families
13	load	TMA producer for input and statistic tiles
14	relay	convert peer DSMEM completion into a local MMA-consumable event
15	empty / register donor	no principal dataflow role

The source uses role-specific setmaxnreg targets. For the audited head-dimension-128 paths:

Path	reduction	compute	load	MMA	empty	Source normalized budget
1CTA	152	136	88	88	24	152 + 2*136 + max(88,88) = 512
2CTA	136	136	104	104	24	136 + 2*136 + max(104,104) = 512

These are source allocation targets, not compiler-reported registers per thread or an occupancy proof. setmaxnreg changes which roles can hold live fragments; the compiled resource record still decides residency.

Repository-authored reconstruction from FA4 Sections 3.2.1-3.2.3 and beta23 source coordinates. TMEM widths are source-derived specialization coordinates, not measured allocation occupancy.

TMEM Offsets Are A Lifetime Schedule

For the audited 1CTA d=128 specialization, the source assigns a 512-column TMEM allocation as follows:

Columns	Serial occupants	Completion edge required before reuse
[0,128)	S -> P	score MMA completes before probability work consumes/reuses it
[128,256)	dV	dV MMA and epilogue consumption complete
[256,384)	dP -> dS -> dQ	dP retires, then dS feeds both gradient MMAs, then dQ may accumulate
[384,512)	dK	dK MMA and epilogue consumption complete

The equality of offsets is not itself proof that aliasing is safe. The proof requires disjoint live intervals and completion edges. In particular, the 1CTA path serializes three occupants in [256,384):

dP completion
    -> elementwise dS formation
    -> dS publication to dK and dQ consumers
    -> dQ MMA may reuse the TMEM columns

The audited 2CTA d=128 map differs:

Effective per-CTA columns	Occupant
[0,64)	current S/P CTA-group accumulator partition
[64,128)	previous-tile dQ accumulator partition
[128,256)	dV
[256,384)	dP -> dS
[384,512)	dK

The source offsets alone do not prove the 64-column physical widths. That reconstruction also consumes the official tcgen05 CTA-group-2 rule: the CTA pair cooperates on one MMA and each CTA receives its accumulator partition. This tighter packing allows previous-tile dQ/dK work to overlap current-tile S/P/dP/dS/dV work in ways the 1CTA alias map cannot express.

There is also a d=192/128 source specialization with different layouts and stage choices. It is a distinct contract, not evidence that the d=128 map generalizes to every head dimension.

One-CTA Prologue, Main Loop, And Tail

At the dependency level, one key tile follows this sequence:

def backward_tile_1cta(q, do, k, v, lse, delta):
    S = mma_qk(q, k)                 # TMEM [0, 128)
    wait_mma(S)
    P = exp_from_lse(S, lse)         # elementwise consumer

    dP = mma_do_vt(do, v)            # TMEM [256, 384)
    dV = mma_pt_do(P, do)            # TMEM [128, 256)
    wait_mma(dP)
    dS = P * (dP - delta[:, None])

    publish_dS_for_gradient_mmas(dS)
    dK = mma_dst_q(dS, q)            # TMEM [384, 512)
    wait_until_dS_operand_retired()
    dQ = mma_ds_k(dS, k)             # reuses [256, 384)

    reduce_dq_tmem_to_global(dQ)
    epilogue_dk_dv(dK, dV)

The real source pipelines multiple tiles. Its prologue creates the first score/probability and gradient states, the main loop overlaps current and previous tiles under separate pipeline phases, and the tail drains outstanding dQ, dK, and dV consumers before storage can be recycled. The pseudocode above preserves ownership and reuse edges but is not a transcription of generated control flow.

Why Two CTAs Must Exchange Half Of dS

Suppose a CTA pair partitions the key dimension. CTA 0 owns columns 0:N; CTA 1 owns columns N:2N. Before dQ, each has

CTA 0: dS[:, 0:N]   and K[0:N, :]
CTA 1: dS[:, N:2N]  and K[N:2N, :]

But every output row requires the complete reduction dimension:

$$dQ=dS[:,0:2N]K[0:2N,:].$$

Leaving the data column-partitioned would give each CTA only a partial dQ, followed by two global atomic contributions for the same output rows. FA4 instead repartitions by output rows:

CTA 0 keeps dS[0:M/2, 0:N]
      receives dS[0:M/2, N:2N]
      computes dQ[0:M/2, :] from the full 2N reduction axis

CTA 1 receives dS[M/2:M, 0:N]
      keeps dS[M/2:M, N:2N]
      computes dQ[M/2:M, :] from the full 2N reduction axis

Each CTA sends one (M/2) x N half. Across the pair, the raw exchange payload is

$$B_{\mathrm{xchg,pair}} =2\left(\frac{M}{2}N\right)b_{dS} =MNb_{dS}.$$

This byte identity is not the paper’s 384-cycle DSMEM entry. The latter is a normalized throughput ledger that includes the paper’s machine-rate assumptions.

Source: Zadouri et al., FlashAttention-4, Figure 3. The paper figure explains the algorithmic repartition; the next diagram reconstructs the pinned source transaction and reduction ownership.

Repository-authored algebra and source trace. DSMEM repairs a partition-axis conflict; the semaphore shown later repairs a floating-point ordering conflict.

The DSMEM Copy Is A Transaction Protocol

The CUDA Programming Guide’s DSMEM contract permits cluster CTAs to access one another’s shared-memory partitions only while the participating blocks exist. The PTX cp.async.bulk contract ties a remote cluster copy to an mbarrier transaction byte count.

The beta23 source implements a multi-role handoff:

compute warpgroup:
    write the exchange half into local sdS_xchg
    elected thread arrives-and-expects peer transaction bytes
    issue cp.async.bulk local SMEM -> remote cluster SMEM

relay warp:
    wait until the remote dS transaction completes
    arrive on a local leader barrier

MMA warp:
    wait for the leader event
    issue dQ MMA only after the peer half is visible

The event order is mechanically checked as

write_xchg
-> expect_tx
-> remote_copy
-> relay_wait
-> relay_arrive
-> mma_wait
-> mma_dQ
-> phase_toggle

Omitting expect_tx loses completion accounting. Omitting the relay edge creates a visibility race. Letting only one CTA advance the phase makes the pair disagree about which generation is live. The source also rejects sparse pairings that do not map both CTAs to the same sparse K/V column; a mismatched pair can wait forever on a peer that is working on another logical tile.

Reconstructing The Two-CTA Traffic Trade

The paper normalizes the 2CTA pair against the same 2560-cycle MMA budget:

Resource	1CTA	2CTA	What changed
MMA compute	2560	2560	same normalized contraction work
SMEM MMA operands	2048	1536	paired MMA shares more operand traffic
dS write	256	256	still required as an MMA operand
dS DSMEM exchange	0	384	new peer repartition cost
dQ write plus read	1024	512	each CTA owns half the output rows
total SMEM/DSMEM ledger	3328	2688	smaller remaining traffic bottleneck
exponential unit	1024	1024	unchanged normalized work

The total shared-memory-side estimate falls by

$$\frac{3328-2688}{3328}=19.23\%.$$

More importantly, the excess above the 2560-cycle MMA budget falls from 768 to 128 cycles:

$$\frac{768-128}{768}=83.33\%.$$

The pair did not remove communication. It exchanged a specific dS half so that it could avoid duplicated operand movement and halve the FP32 dQ staging/global-reduction surface. Normalizing one pair over two N partitions, the oracle counts

1CTA plan: 65,536 dQ atomic elements
2CTA plan: 32,768 dQ atomic elements
ratio:      0.5

This is a count of reduction elements under the paper tile mapping, not a claim that atomic latency or end-to-end backward time is exactly halved.

dQ Reduction: TMEM To RMEM To SMEM To Global

The output of the dQ MMA is not directly a final global tensor. In beta23:

reduction warps 0-3
    load CTA-owned dQ rows from TMEM into registers
    repack and stage them into shared memory
    execute a shared-memory fence
    synchronize the 128 reduction threads
    issue bulk global FP32 reduce-add
    wait for completion before stage reuse or semaphore release

The shared fence and barrier are correctness edges. Performing the global add before all four reduction warps have published their slice can expose partially initialized shared memory. Reusing the stage before the bulk operation completes can overwrite bytes still in flight.

Determinism Requires A Writer Order

Global gradient accumulation combines contributions from multiple tiles or query heads. Floating-point addition is not associative. The oracle’s four FP32 contributions

[100000000.0, 1.0, -100000000.0, 3.0]

sum to 4.0 in one arrival order and 3.0 in the enforced writer order. A deterministic kernel therefore needs more than a deterministic mathematical expression; it needs a deterministic physical reduction sequence.

The pinned source follows this protocol:

wait_eq(semaphore, my_lock_value)
stage_dq_tmem_to_smem()
shared_fence_and_reduction_warp_barrier()
bulk_global_reduce_add_f32()
wait_for_bulk_reduce_completion()
arrive_inc(semaphore)

The order of those calls matters. Releasing the semaphore before the global reduce-add completes allows the next writer to overlap the same accumulator in an order that is no longer the declared sequence.

The source uses different lock coordinates for different ownership patterns:

Case	Ordered coordinate	Source-derived lock idea
dense dQ	K/V tile ordinal	lock = n_block
causal SPT dQ	valid K/V tiles in reverse writer order	lock = n_block_max_for_m_block - 1 - n_block
GQA dK/dV	query head within a shared K/V head	lock = head_idx % qhead_per_kvhead

For the causal example with valid physical writers [1,2,3], reverse scheduling produces execution order [3,2,1] and lock values {3:0, 2:1, 1:2}. Sparse or truncated schedules create another hazard: if a physical writer slot is absent, later writers cannot wait for an increment that will never occur. Beta23 contains tail bridging that advances missing slots. The oracle intentionally removes that bridge and observes a deadlock after the first writer.

The deterministic path pays for waiting, fences, and serialized add order. The nondeterministic path lets arrival order select the FP32 add order and can be faster. They are different schedules and must be benchmarked as such.

Variants And Design Alternatives

Design axis	Alternative	Benefit	Cost or constraint
CTA group	1CTA	simpler local ownership and no DSMEM exchange	duplicated traffic and more dQ reductions
CTA group	2CTA	lower normalized SMEM ledger and half the dQ atomic elements	cluster admission, peer liveness, DSMEM transaction, stricter phases
head dimension	d=128 map	audited compact TMEM layout	not universal
head dimension	d=192/128 source specialization	supports another shape family	different layouts/stages; separate proof required
reduction order	nondeterministic bulk add	avoids semaphore serialization	run-to-run add order may vary
reduction order	deterministic semaphore	repeatable declared writer order	wait/fence/tail-bridge overhead
head sharing	MHA	one query head per K/V head	less cross-head reduction sharing
head sharing	GQA/MQA	K/V gradients combine query-head contributions	additional deterministic head-order coordinate
sparsity	dense/causal	regular or monotone writer sets	causal order still needs valid-slot accounting
sparsity	block sparse / variable length	skips work	paired CTAs and semaphore slots must agree on the same physical work set

The decisive design choice is not “use 2CTA because it is newer.” It is whether the saved local traffic exceeds the cost of cluster admission, exchange, synchronization, and any occupancy or tail loss for the actual shape.

Executable Reconciliation

The backward oracle reconstructs both the algebra and the ownership schedule on deterministic CPU fixtures:

Check	Result
one-CTA paper roofline maximum absolute error	0 cycles
tiled dK versus direct reference	5.55e-17 max absolute error
tiled dQ versus direct reference	1.11e-16 max absolute error
tiled dV versus direct reference	1.11e-16 max absolute error
exact 2CTA row/column repack	0 max absolute error
omit peer dS exchange	0.6799 max absolute dQ error
exchange the wrong row half	1.2337 max absolute dQ error
swap peer K column order	1.4536 max absolute dQ error
valid relay event sequence	accepted
omit expected bytes / relay arrival / MMA wait	rejected
valid TMEM live intervals	no concurrent overlap
force concurrent 1CTA dP/dS/dQ alias	rejected
deterministic dense and causal writer orders	completed
missing writer slot without tail bridge	deadlock detected

These results are stronger than prose examples and weaker than a compiled kernel run. The oracle does not lower CuTe DSL, model TMEM latency, prove compiler scheduling, or profile B200 hardware.

Paper, Source, And Runtime Boundaries

Claim	Highest evidence reached here
five-MMA and softmax-gradient algebra	derivation plus deterministic oracle
paper cycle ledger	exact arithmetic reproduction
warp IDs, register targets, TMEM offsets, relay, and reduction calls	pinned beta23 source
CTA-group accumulator partition and DSMEM lifetime	official NVIDIA documentation plus source reconstruction
legal generated barriers for every specialization	not established
compiled register count, occupancy, selected kernel, overlap, and timing	not measured

Renderer-safe formula ledger:

S = alpha * Q @ K^T
P = row_softmax(S)
O = P @ V

dP = dO @ V^T
delta_i = sum_j P_ij * dP_ij = dot(dO_i, O_i)
dS = P * (dP - delta[:, None])
dV = P^T @ dO
dK = alpha * dS^T @ Q
dQ = alpha * dS @ K

T_MMA = 10 * M * N * d / 8192
T_SMEM_MMA = (4*M*d + 3*N*d + M*N) / 64
T_dS_write = M*N / 64
T_dQ_write_read = M*d / 16
T_exp = M*N / 16

B_xchg_pair = M * N * bytes_dS
T_tile_lower_bound ~= max(T_MMA, T_SMEM_total, T_exp)

Backward Failure Lab

Failure	Why a plausible implementation can still be wrong	Required proof
count only four MMAs	forgets score recomputation because forward did not retain S	logical graph and FLOP derivation
materialize P globally	preserves algebra but abandons FlashAttention’s state contract	tile-local reconstruction trace
reuse equal TMEM offsets concurrently	coordinates match but lifetimes overlap	interval plus completion-edge proof
infer 2CTA widths from offsets alone	ignores CTA-group accumulator partition	official CTA-group contract plus source layout
omit peer half of dS	each CTA computes only a partial reduction axis	exact repack oracle
exchange the same half twice	byte counts and shapes still look valid	row/column ownership check
issue dQ MMA before relay arrival	peer data may not be visible	expect-copy-relay-MMA event trace
pair different sparse K/V columns	each CTA waits on the wrong peer generation	pair-descriptor admission
reduce before shared fence/barrier	global bulk add sees incomplete staging	reduction-warp publication proof
release lock before bulk completion	next writer changes the declared FP32 order	wait-before-arrive protocol
leave a missing semaphore slot	a later writer waits forever	tail bridge or compact writer numbering
call the CPU oracle B200 evidence	arithmetic checks pass without a compiled kernel	artifact plus profiler evidence

Backward Boundary And Release Gate

Publish this loop only if:

• the softmax Jacobian produces the final dS formula and zero-row-sum invariant;
• delta = row_sum(P*dP) = row_dot(dO,O) is derived and tested;
• all five MMAs have tensor shapes, producers, consumers, and precision boundaries;
• the one-CTA roofline is derived from operand traffic rather than copied as a table;
• paper lower bounds remain distinct from profiler time;
• beta23 warp roles and source register targets are not called compiled occupancy;
• TMEM reuse is justified by lifetimes and completion edges;
• 1CTA, 2CTA d=128, and the separate d=192/128 specialization are not conflated;
• the half-dS exchange is derived from the reduction-axis conflict;
• DSMEM transaction bytes, relay completion, and MMA consumption are ordered;
• dQ staging includes TMEM, RMEM, SMEM, fence/barrier, and global reduce-add;
• deterministic dense, causal SPT, sparse-tail, and GQA orders are distinguished;
• all negative mutations fail mechanically;
• the CPU oracle is not presented as a Blackwell artifact or timing result;
• persistent scheduling, CuTe artifact integration, and benchmarks remain deferred to 0x0B-0x0D.

0x0B Persistent Scheduling, LPT, And Variable-Work Geometry

Owner question: once the inner kernel is pipelined, how should a finite set of unequal attention tiles be ordered and assigned so that the last SM does not inherit all of the long work?

This section owns global tile-cost estimation, LPT coordinate order, persistent work progression, cache-aware head/batch swizzles, CLC work assignment, and paper-to-release scheduler drift. It consumes the per-tile pipelines derived in 0x03-0x0A. It does not rederive 0x0A’s dQ semaphore or shortest-processing-time-first reduction protocol, and it does not reopen M4.5’s framework-level backend admission.

Four Terms That Sound Similar But Are Not

Term	What it controls	State that changes	What it does not prove
Persistent CTA scheduling	how one resident CTA obtains more than one logical tile	current linear tile/work descriptor	that the tile order is load-balanced
LPT ordering	which estimated-cost tile is made available first	logical coordinate linearization	that CTAs remain resident or steal work
CLC dynamic scheduling	how hardware-assisted work descriptors are requested and returned	producer/consumer work token	that the static L2 swizzle is preserved
Deterministic SPT reduction	which backward writer reaches a shared dQ row first	semaphore turn and reduction order	that forward or backward work is globally longest-first

The distinction matters in the pinned source. StaticPersistentTileScheduler caps the physical grid by the SM count and advances a resident CTA by grid stride. SingleTileLPTScheduler maps a linear tile index through an L2 swizzle and optional block reversal; its static path performs one logical tile per CTA, while its CLC path consumes hardware-assigned work. One class name or one is_persistent flag therefore cannot describe the whole scheduling contract.

Deriving A Causal Tile-Cost Proxy

Let a query tile contain rows

$$q \in [mM,\min((m+1)M,S_q)),$$

and let each K/V tile contain $N$ key positions. For aligned causal self-attention without a local window, the largest key index visible to query tile $m$ is approximately the last query position in the tile. The number of K/V tiles that can contain a valid score is therefore

$$J_m = \min\!\left( \left\lceil\frac{S_k}{N}\right\rceil, \left\lceil\frac{\min((m+1)M,S_q)}{N}\right\rceil \right).$$

For $M=N$ and a full square sequence, this reduces to

$$J_m=m+1.$$

The real cycle cost contains a fixed tile prologue/epilogue plus per-K/V-tile loads, two MMAs, mask work, softmax work, and possible sparse or tail corrections:

$$w_{b,h,m} \approx c_{\text{fixed}} + J_{b,h,m} \left( c_{QK}+c_{\text{softmax}}+c_{PV}+c_{KV\text{-load}} \right).$$

J_m is useful because it is cheap metadata and preserves the dominant monotone ordering for the causal dense case. It is not a universal cycle model. A sliding window clips $J_m$; block sparsity replaces it with selected-block count; variable lengths make it batch-dependent; SplitKV and two-CTA modes change the cost per selected block.

If $P$ identical worker slots receive tile set $\mathcal{T}$, the makespan objective is

$$C_{\max} = \max_{p\in\{1,\dots,P\}} \sum_{i\in A_p} w_i,$$

with the unavoidable lower bound

$$C_{\max}^{*} \ge \max\!\left( \frac{\sum_i w_i}{P}, \max_i w_i \right).$$

This is why utilization at the beginning of a launch is not enough. The kernel completes when the most-loaded worker finishes.

Worked Makespan: Lexical Order Versus LPT

For six causal query blocks with cost proxy [1,2,3,4,5,6] and three worker slots, a first-free list scheduler gives:

Order	Worker loads	Makespan	Interpretation
lexical m=0..5	5, 7, 9	9	the late wave contains the longest tiles
LPT m=5..0	7, 7, 7	7	long work starts early; short tiles fill the tail

The lower bound is max(21/3, 6) = 7, so LPT reaches it in this toy. That is an accounting example, not the source of the paper’s measured speedups. Its purpose is to expose the tail mechanism that a TFLOPS bar hides.

def list_schedule(tile_costs, logical_order, worker_count):
    loads = [0] * worker_count
    schedule = []
    for tile in logical_order:
        worker = min(range(worker_count), key=lambda p: (loads[p], p))
        start = loads[worker]
        finish = start + tile_costs[tile]
        loads[worker] = finish
        schedule.append((tile, worker, start, finish))
    return schedule, max(loads)

lexical = range(6)              # costs 1,2,3,4,5,6
lpt = reversed(range(6))        # costs 6,5,4,3,2,1

Why Naive Global LPT Can Still Lose

The FA4 paper’s Section 3.3 does not merely reverse a flat list. A global longest-first sort can interleave batches whose K/V tensors are unrelated or bring too many K/V heads into the active set. It may reduce arithmetic tail length while increasing L2 misses.

For one batch and one K/V head, a rough resident K/V working-set estimate is

$$W_{KV,\text{head}} \approx S_k(d_{qk}+d_v)b,$$

so a head section should satisfy

$$H_{\text{section}}W_{KV,\text{head}} \lesssim C_{L2,\text{budget}}.$$

The paper uses this hierarchy:

for batch in batches:                       # unrelated K/V batches stay outside
    for head_section in L2_fitting_sections:
        for mblock in reverse(query_blocks): # longest causal work first
            for head in head_section:
                issue(batch, head, mblock)

The beta23 fixed-length source mechanically reconstructs the same shape of static coordinate map: it flattens batch/head sections, chooses a power-of-two head swizzle, reverses mblock when lpt=True, and treats the residual head section separately. The source sets its numerical L2 heuristic to 50 * 1024 * 1024 bytes even though one adjacent comment still says 40 MB; that constant is a release-specific heuristic, not a portable hardware capacity statement.

For MQA/GQA, the paper adds a reuse constraint: traverse all query heads sharing one K/V head before moving the K/V-head work coordinate. This reduces repeated K/V working-set turnover. It is a cache-locality choice layered on top of the cost order, not a change to attention semantics.

The paper reports, on H200 with BF16 and head dimension 128, 4-8% FLOPS gain for MHA and 7-14% for MQA8 from the described scheduling order. Those values are source-reported scheduling measurements on the paper’s configurations, not expected gains for beta23, B200, arbitrary sequence distributions, or end-to-end serving.

Static Persistence And CLC Work Stealing

The pinned source implements at least two different ways to move from one logical work tile to another.

Static persistent grid-stride path. StaticPersistentTileScheduler.get_grid_shape caps grid x at the number of available SM slots, rounded for cluster shape. Each resident CTA then advances:

next_linear_tile = current_linear_tile + physical_grid_x

This reduces the physical launch grid and lets each resident CTA consume several tiles. The assignment is deterministic from grid geometry, but it can still suffer if every stride column receives a bad mixture of long and short tiles.

CLC dynamic path. SingleTileLPTScheduler and SingleTileVarlenScheduler also accept SchedulingMode.CLC. Their protocol is:

work = clc.initial_work_tile_info()
while work.is_valid:
    consume(map_hardware_work_to_attention_coords(work))
    clc.prefetch_next_work()
    clc.consumer_wait()
    work = clc.get_current_work()
    clc.consumer_release()
clc.producer_tail()

The CLC response contains raw work-grid coordinates. The beta23 source explicitly says that this path does not apply the static L2 swizzle because hardware decides work order; it only applies cluster conversion, optional LPT block reversal, and SplitKV unpacking. Dynamic assignment can improve tail balance, but it can also sacrifice the cache order that made static LPT useful.

That tradeoff appears in the host selector. In beta23, a requested CLC scheduler is disabled for variable-length MHA and dense noncausal attention because source comments record regressions from additional K/V blocks in flight, worse L2 behavior, or pure work-stealing overhead. use_clc_scheduler is then included in the forward compile key, so changing this policy changes artifact identity rather than only runtime metadata.

Variable Lengths: Paper Algorithm And beta23 Source Diverge

For sequence lengths [128,768,256,512] with M=128, the query worktile counts are [1,6,2,4]. A paper-style LPT preprocessing pass produces a virtual-to-actual map:

virtual batch:  0  1  2  3
actual batch:   1  3  2  0
worktiles:      6  4  2  1

The paper proposes a preprocessing kernel that sorts batches by estimated maximum per-worktile execution time, writes that map, and allows the metadata to be cached. That design is especially relevant to mixed prefill/decode batches, where lexical request order has no reason to correlate with cost.

The pinned beta23 SingleTileVarlenScheduler should not be described as that exact implementation. Its main coordinate map reads cu_seqlens_q or seqused_q, computes per-batch query-block counts with warp prefix sums, scans groups of 31 batches to locate the current flat tile, and then applies per-batch head swizzle and optional block reversal. The main scheduler does not consume the paper’s virtual-to-actual sort map. A varlen_batch_idx appears in the separate forward-combine surface, but that is not evidence that the main attention scheduler implements the Section 3.3 preprocessing contract.

This is a useful form of source drift:

Contract	Batch order	Cost order inside batch	Evidence
FA4 paper varlen design	sorted virtual batches	longest estimated tile first	paper algorithm
beta23 main varlen scheduler	original metadata order located by prefix sums	optional reverse block and head swizzle	pinned source
deployment	whichever compiled path actually runs	profiler-observed	not measured here

Deterministic Reduction Is A Consumed Boundary

Section 0x0A already derives why deterministic dQ uses semaphores, fences, a declared writer order, causal SPT, sparse-tail bridging, and GQA-aware turns. This section consumes only the consequence:

LPT:
    tries to reduce global makespan by starting expensive work early

deterministic causal SPT:
    tries to make the first dQ writer arrive without waiting and fixes FP32 order

The two objectives can conflict. The paper’s deterministic schedule combines K/V-block order, query-block order, batch/head swizzle, and reduction turn order rather than applying one universal sort.

Source: Zadouri et al., FlashAttention-4, Figure 8, cropped from the CC BY-NC-SA 4.0 paper. Left: noncausal batch/head swizzle versus naive order. Right: causal SPT, reverse-mblock LPT, naive LPT, and naive order. Values are source-reported B200 FP16/BF16 deterministic-backward throughput for head dimension 128; they are not beta23 measurements.

For mobile-readable spot checks, the figure reports:

Sequence	Noncausal swizzle / naive	Causal SPT / reverse-LPT / LPT / naive
1k	630 / 500	470 / 450 / 440 / 250 TFLOP/s
8k	850 / 730	920 / 790 / 720 / 350 TFLOP/s
32k	880 / 880	950 / 830 / 760 / 540 TFLOP/s

The convergence of noncausal swizzle and naive order at 32k is as important as the 1k gap: scheduler value depends on wave count, cache behavior, and tail geometry.

Scheduling Failure Closure

Mutation	What still looks plausible	What fails
launch causal mblock in increasing order	all tiles are correct	late workers inherit the longest mainloops
globally sort all heads and batches by cost	arithmetic tail may shrink	unrelated K/V working sets thrash L2
call grid-stride persistence LPT	CTAs process many tiles	logical order can remain poorly balanced
call CLC work stealing the static L2 swizzle	both return legal coordinates	cache locality contract is different
call beta23 varlen prefix sums the paper sort	batch coordinates are recovered	no virtual-to-actual LPT ordering exists
apply forward LPT to deterministic dQ turns	long tiles start early	first writer can stall and FP32 order changes

0x0C CuTe DSL, Dispatch, And The Artifact Boundary

Owner question: which facts are proven by Python source, which are fixed by a compile key, and which require the generated artifact and the runtime launch?

CuTe DSL is Python-hosted metaprogramming, not interpreted device execution. NVIDIA’s CuTe DSL code-generation documentation describes Python execution and tracing into compiler IR, progressive lowering, translation to PTX/SASS, device-binary construction, module load, and launch. The debugging contract exposes retained IR, PTX, CUBIN, SASS, and Python-to-device line information.

Source-To-Execution Proof Ladder

Python API call
    -> host validation and dispatch predicates
    -> compile key and specialization cache
    -> CuTe IR / MLIR lowering
    -> PTX
    -> ptxas -> CUBIN / SASS
    -> JIT executor and loaded kernel symbol
    -> profiler trace + numerical comparison + timing distribution

Each arrow can fail independently. An API may admit a shape that fails compilation. A compiled executable may not be the one selected after a cache-key bug. A selected kernel may run correctly but spill or occupy fewer CTAs than the source target. A fast kernel-only timing may still lose end-to-end after compilation, metadata preprocessing, or fallback.

Paper-Era Hardware Matrix

The FlashAttention-4 paper targets B200/GB200 Blackwell. FlashAttention-3’s paper contract targets Hopper, while FlashAttention-4’s algorithm and reported resource model target Blackwell. Those paper coordinates establish algorithmic intent and source-reported measurements; they do not establish the dispatch surface of a later package release.

Paper coordinate	Hardware coordinate	Contract used here
FlashAttention-3	H100/H800-class Hopper	TMA/WGMMA pipeline, warp specialization, FP8 method, and paper measurements
FlashAttention-4	B200/GB200-class Blackwell	TMEM/tcgen05 pipeline, two-CTA schedule, software exponential path, and paper measurements

Pinned Artifact Scope

The pinned beta23 package exposes SM90, SM100, and SM120-family paths as well as a separate SM80 forward path. The broader beta23 source surface does not retroactively change the FA4 paper’s hardware contract. Conversely, the paper does not prove that every feature in beta23 uses the paper’s exact schedule: that requires source-path selection, generated-artifact inspection, and runtime evidence for the chosen specialization.

The Release Pin Moved During This Audit

The implementation anchor remains pre-release fa4-v4.0.0.beta23, published 2026-07-22 at commit 2409214a03797b168f648ea30df1adbc09ce658a. As of this audit on 2026-07-30, the official repository has already published fa4-v4.0.0.beta24, dated 2026-07-29 at commit 849f660f73b176e5ad5670e7f822c7fa9f3eaf8b.

Layer	Immutable identity	What this article may claim
FA4 paper	arXiv 2603.05451v1, audited PDF SHA-256 7c053b...f10eb	paper algorithms and source-reported measurements
implementation anchor	beta23, commit 2409214a...	source predicates and source structure extracted below
drift head	beta24, commit 849f660f...	release-note drift only, not beta23 behavior
local execution	none	no compiled kernel, artifact, profile, or timing claim

beta24’s official release notes include variable-length backward score-mod support, combine-kernel scheduling changes for CUDA Graph capture, and FP8 numerical changes. The beta23 and beta24 tile_scheduler.py blob is unchanged, but interface.py changed. That is enough to keep the matrix pinned: release drift can alter validation, compile keys, numerical paths, or wrapper integration even when one scheduler file is identical.

beta23 Interface Admission Is A Predicate Set

The beta23 forward interface first validates tensor shapes, dtype equality, integer metadata layout, CUDA placement, architecture family, head ratio, and optional-state compatibility. It then derives tile sizes, SplitKV, two-CTA eligibility, sparse normalization, scheduler mode, and finally a compile key.

The source-admitted surface is broader than the FA4 paper:

Surface	beta23 source predicate	Important exclusions
forward architecture	arch // 10 in [8,9,10,11,12]	each family selects a different class and feature set
backward architecture	arch // 10 in [9,10,11,12]	SM80 forward admission does not imply beta23 backward admission
input dtype	FP16, BF16, FP8 E4M3FN/E5M2 at API boundary	FP8 is forward-only and restricted to SM10x source path
SM90 head dimensions	aligned 8..256 for Q/K and V	profitable specialization is not guaranteed
SM100/110 head dimensions	aligned 8..128, (192,128), selected 64/512, or (256,256)	special paths add feature exclusions
SplitKV forward	SM10x/11x conditional	source rejects it on SM80, SM90, and SM120; shape can force one split
paged K/V forward	SM90 and SM10x/11x conditional	rejected on SM80/SM120; page-size/path constraints remain
Blackwell 2-CTA forward	noncausal, nonlocal, fixed-Q, nonsplit, dense, compatible shapes	host/env flags and dedicated d=256 rules can override
static persistent forward	SM10x/11x noncausal, nonlocal, fixed-Q, nonsplit	persistence does not imply CLC or LPT
deterministic backward	selected SM90/SM10x/SM11x paths	explicitly rejected on SM120; sparse mode needs write-order metadata

This is a source-admission matrix. It does not mean every admitted combination was compiled in this audit, passes correctness tests, selects the paper schedule, or is faster than another backend.

Why The Compile Key Is Part Of Correctness

The beta23 forward key has dozens of fields. Listing them flat hides why they exist, so group them by the generated code they can change:

numeric and shape:
    dtype, d_qk, d_v, qheads_per_kvhead, architecture

semantics:
    causal, score_mod hash, mask_mod hash,
    block-sparse presence/broadcast pattern, auxiliary captures

optional tensor state:
    LSE presence, cu_seqlens/seqused presence, page table,
    window sides, learnable sink, FP8 descales, sparse offsets

tile and schedule:
    tile_m, tile_n, q_stage, thread count, SplitKV, pack-GQA,
    paged non-TMA mode, 2CTA, subtile factors,
    PV RS/SS placement, intra-WG overlap, CLC

specialized paths:
    Q/QV presence, top-k gather length, sparse-KV flags, log level

The operational pattern is:

compile_key = derive_constexpr_identity(validated_call, selected_policy)

if compile_key not in compile_cache:
    kernel_object = select_architecture_and_feature_class(compile_key)
    compile_cache[compile_key] = cute.compile(
        kernel_object,
        fake_or_specialization_tensors,
        options="--enable-tvm-ffi",
    )

executor = compile_cache[compile_key]
executor(runtime_tensors_and_scalars)

If a value changes control flow, instruction family, static tensor layout, optional pointer ABI, or synchronization topology, it belongs either in the key or in a validated runtime contract. Omitting a compile-time branch can return a stale executor whose Python call signature remains shape-compatible while its device assumptions are wrong.

The reverse error is also costly: adding values that do not affect generated code can explode cache cardinality and first-use compile latency. Artifact integration is therefore a partitioning problem between:

constexpr identity -> one generated executor
runtime metadata   -> many legal calls to that executor

The Compiler Artifact Must Be Retainable

The paper describes:

CuTe DSL Python -> PTX -> ptxas -> SASS

Current official CUTLASS documentation makes the proof boundary more explicit. A compiled JIT executor can expose generated MLIR/PTX/CUBIN, and debug options can retain IR, PTX, CUBIN, and SASS. A release-quality artifact manifest should therefore contain:

source:
    repository, tag, commit, source-file blob SHAs

compiler:
    nvidia-cutlass-dsl, CUDA/ptxas, target SM, options, environment

specialization:
    canonical compile key, input ABI, selected class and scheduler

artifact:
    IR/PTX/CUBIN hashes, loaded module and kernel symbol, SASS hash

execution:
    driver/GPU identity, resource attributes, profiler trace, correctness result

beta23’s package pin is itself multi-component: Python >=3.10, nvidia-cutlass-dsl==4.6.0.dev0, apache-tvm-ffi>=0.1.12,<0.2, and quack-kernels>=0.5.3. The unbounded upper range on Quack means the FA4 tag alone does not determine every helper implementation. A strict reproduction should lock the resolved environment, not just copy the tag.

Paper, Source, Artifact, And Runtime Claim Matrix

Claim	Paper	beta23 source	Compiled artifact	Runtime/profile
LPT design exists	proves	contains related mappings	not required	measures effect
exact scheduler variant ran	no	eligible only	symbol/code can prove	profiler confirms
shape is accepted	paper cases only	host predicate proves	compile may still fail	execution proves
WGMMA/tcgen05 instruction emitted	intent/algorithm	source intent	PTX/SASS proves	counters corroborate
register/SMEM/TMEM residency	model/targets	annotations/heuristics	ptxas metadata partly proves	profiler attributes prove
beta23 reaches paper TFLOPS	no	no	no	requires beta23 timing manifest

The evidence produced for this article reaches the paper and source columns. It deliberately does not claim local compilation or execution.

Release And Re-Audit Policy

A current-support statement expires when any of these changes:

attention release or source commit
compiler/CUDA/ptxas version
target architecture
compile-key schema or selected specialization
dependency resolution
benchmark numerator or timing boundary

The re-audit sequence is mechanical: extract the new interface predicates, regenerate the support matrix, diff source blob SHAs and compile-key groups, retain a representative artifact, rerun correctness, then benchmark the selected symbol. A release note is a drift signal, not a substitute for that sequence.

0x0D How To Read FA3 And FA4 Benchmarks

Benchmark Numerators, Timing Boundaries, And Reconciliation

Owner question: when a graph says TFLOP/s or 1.3x, which operations were counted, which wall-clock interval was timed, and which exact artifact and baseline own the number?

Derive The Counted FLOPs Before Reading The Bar

For batch $B$, query length $S_q$, key length $S_k$, $H$ heads, Q/K head dimension $d_{qk}$, and V dimension $d_v$, dense forward attention contains two dominant matmuls:

$$F_{QK^\top}=2BHS_qS_kd_{qk},$$ $$F_{PV}=2BHS_qS_kd_v.$$

Therefore

$$F_{\text{fwd,dense}} = 2BHS_qS_k(d_{qk}+d_v).$$

When $d_{qk}=d_v=d$ and $S_q=S_k=S$, this becomes the paper’s familiar per-batch form with the restored batch multiplier:

$$F_{\text{fwd,dense}}=4BHS^2d.$$

For exact aligned causal self-attention, the valid score-pair count is triangular:

$$N_{\text{pairs,causal}}=\frac{S(S+1)}{2},$$

so an exact useful-work numerator would replace $S^2$ by $S(S+1)/2$. The FA4 paper uses the simpler convention of dividing dense FLOPs by two. At long sequence length the relative difference is $1/S$; at short length or with tile padding it can be visible.

The paper counts backward as 2.5 * forward because the backward algorithm performs five dominant matmuls while forward performs two:

$$F_{\text{bwd,paper}}=2.5F_{\text{fwd,paper}}.$$

This omits non-matmul instructions, preprocessing, reduction traffic, and recomputation details from the numerator even though they remain in elapsed time. TFLOP/s is therefore an accounting throughput, not a statement that every cycle issued tensor-core FLOPs.

Useful, Scheduled, And Dense-Equivalent Numerators

One number is not enough for local windows, sparsity, ragged batches, or padded tiles:

Numerator	Counts	Useful for	Failure if mislabeled
useful FLOPs	semantically valid score/output pairs	algorithmic efficiency	hides executed padding or masked work
scheduled FLOPs	worktiles/instructions actually issued under a declared model	kernel efficiency	requires accurate path-specific accounting
dense-equivalent FLOPs	dense reference work regardless of sparsity	effective application rate	can exceed hardware peak without being physical FLOPs

For a block-sparse kernel, dense-equivalent TFLOP/s can legitimately exceed a hardware dense peak because the numerator credits skipped work. Calling it measured tensor-core throughput would be wrong.

A Numeric Reconciliation

Take B=32, S_q=S_k=1024, H=16, d=128. The counted dense forward work is

4 * 32 * 1024 * 1024 * 16 * 128
= 274,877,906,944 FLOPs

Under the paper’s causal-half convention:

Mode	Counted FLOPs	Example time	Derived accounting throughput
dense forward	274.878e9	0.20 ms	1374.39 TFLOP/s
causal forward	137.439e9	0.20 ms	687.19 TFLOP/s
dense backward	687.195e9	0.45 ms	1527.10 TFLOP/s
causal backward	343.597e9	0.45 ms	763.55 TFLOP/s

These times are synthetic oracle inputs, not measurements. The point is that reusing dense forward FLOPs for the causal bar doubles the reported number, while forgetting the backward 2.5x factor understates it.

Kernel Time Is Not Compile Time Or End-To-End Time

At least four clocks can be relevant:

developer compile time:
    source specialization -> executable artifact

cold first-call latency:
    validation + cache lookup/miss + compile/load + launch

warm kernel time:
    already-compiled selected symbol only

end-to-end attention latency:
    metadata preprocessing + packing/paging + kernel(s) + combine + framework overhead

FA4 Table 4 reports one-kernel developer compile time:

Method	Forward compile	Backward compile	Reconciled speedup
FA3 C++ templates	55 s	45 s	baseline
FA4 CuTe DSL	2.5 s	1.4 s	22.0x, 32.14x (32x rounded in paper)

This supports the paper’s 20-30x iteration claim. It does not imply a faster GPU kernel, a 32x first-request improvement, or a complete package build that is 32x shorter. The release’s benchmark.py helper also has its own timing scopes; its presence does not prove that beta23 recreated Table 4 or the paper’s B200 protocol.

FA3 Source Results Remain A Hopper Contract

The FA3 paper evaluates H100 80GB SXM5 with FP16 inputs across sequence lengths 512 through 16k, fixed total-token configurations, and several head dimensions. It reports:

up to 740 TFLOP/s, about 75% of the cited H100 theoretical rate
about 1.5-2.0x FA2 forward in tested configurations
about 1.5-1.75x FA2 backward in tested configurations
FP8 forward close to 1.2 PFLOP/s in reported cases

Source: Shah et al., FlashAttention-3, Figure 5. These are forward-kernel measurements on H100 80GB SXM5 for the paper’s declared inputs, shapes, and baselines. They are not end-to-end model throughput and do not predict a Blackwell result.

The FA4 paper notes that its compared FA3 implementation does not run on B200. That is a paper-era artifact statement: WGMMA, TMEM, CTA-group instructions, compiler support, and synchronization contracts are architecture-specific. It does not mean that no later FA3-derived or CuTe implementation can target Blackwell.

FA4 Paper Benchmark Manifest

The main text’s benchmark contract is:

Field	Paper value
hardware	B200 in main text and Figures 4-8
input	BF16 in Section 5; captions say FP16/BF16
sequence lengths	1k, 2k, 4k, 8k, 16k, 32k
batch	chosen so total tokens equal 32k
hidden dimension	2048
standard heads	d=64, H=32; d=128, H=16
DeepSeek shape	(d_qk,d_v)=(192,128), H=16
mask	causal and noncausal
forward numerator	4 * S^2 * d * H, with batch implied by problem and causal divided by two
backward numerator	2.5 * forward
compared systems	PyTorch, FA2, Triton, Gluon, cuDNN 9.13

There is an unresolved metadata inconsistency inside the paper. The paper’s main text and plotted labels identify B200, while Appendix A.1 records B100; the full conflict is preserved below before these values are interpreted as a reproduction contract.

It reports up to 1613 TFLOP/s, about 71% of the cited B200 theoretical rate, up to 1.3x cuDNN 9.13, and up to 2.7x Triton. These are best-case, source-reported paper claims.

Source: Zadouri et al., FlashAttention-4, Figure 4. The plots report B200 forward accounting throughput with head dimension 128. The paper states that newer cuDNN versions incorporated related techniques and can approach FA4, so this is a versioned comparison, not a permanent ranking.

Preserve The Paper’s Metadata Contradiction

Appendix A.1 says:

GPU: B100 180GB SXM6 (1000W)
warmup: 5 runs
timed repetitions: 10
statistic: arithmetic mean
CUDA: 13.1
FlashAttention: 2.8.3
Triton: 3.6
PyTorch: 2.10.0
CuTe-DSL: 4.4.1
cuDNN: 9.13 in main comparison; 9.19.1.2 as latest appendix entry

The main text and every relevant figure label say B200, while the appendix says B100 and dates the software snapshot to “March 2025” despite arXiv v1 being submitted in March 2026. This article does not guess which appendix labels are typographical errors. It does not silently replace B100 with B200, infer the intended date, or merge the two manifests. Until an author correction resolves the conflict, the appendix discrepancy remains part of the benchmark record and prevents the appendix from serving as a complete reproduction contract.

Baseline Reconciliation

Speedup is a ratio of two measured intervals:

$$S_{A/B}=\frac{t_B}{t_A}.$$

It is only stable when numerator semantics and timing boundaries match. Before quoting a bar, verify:

same GPU, power mode, and clocks
same B, Sq, Sk, Hq, Hkv, dQK, dV
same dtype, scaling, mask, window, sparsity, and deterministic mode
same warmup, synchronization, repetitions, and statistic
same allocation/preprocessing/compile inclusion boundary
same correctness threshold
exact baseline and tested-kernel versions

The paper itself says newer cuDNN releases adopted related techniques. A comparison against cuDNN 9.13 does not establish the same ratio against 9.19.1.2, beta23, beta24, or a framework backend that selects another kernel.

Source-Reported, Release-Audited, And Locally Measured Must Stay Separate

Statement	Correct label
FA4 reaches up to 1613 TFLOP/s	source-reported paper peak
beta23 accepts this architecture/shape/feature predicate	pinned-source audit
beta23 compiles to this symbol and SASS	requires retained local artifact
that symbol reaches X TFLOP/s here	requires local timing manifest
the deployment gains Y tokens/s	requires end-to-end service measurement

No Blackwell or Hopper GPU artifact was compiled or timed for this loop. The local oracle checks schedule accounting, compile-time ratios, and FLOP numerators only.

Reproduction And Rollback Manifest

For one benchmark row, record:

identity:
    paper/release/commit, source blob SHAs, dependency lock
    IR/PTX/CUBIN/SASS hashes and loaded kernel symbol

environment:
    GPU product, SM count, memory, clocks, power mode
    driver, CUDA, ptxas, CuTe DSL, PyTorch, framework

dispatch:
    requested backend, admitted path, compile key
    scheduler, 1CTA/2CTA, SplitKV, persistent/CLC, fallback

problem:
    B, Sq, Sk, Hq, Hkv, dQK, dV
    dtype/scales, causal/window/sparse, varlen/paged state

measurement:
    FLOP convention, warmup, repetitions, synchronization
    statistic and distribution, compile/preprocess/allocation boundary

correctness:
    reference, tolerance, max/mean error, deterministic behavior

A release candidate should also carry a fallback and rollback result: force an unsupported shape, verify that the declared backend rejects or falls back visibly, and verify that reverting the release restores the previous artifact and performance distribution. Otherwise a deployment regression can masquerade as natural workload variance.

Scheduling, Artifact, And Benchmark Release Gate

• causal tile cost and makespan lower bound are derived;
• lexical, LPT, cache-aware LPT, persistence, CLC, and deterministic SPT are distinct;
• fixed-length and variable-length scheduler coordinate traces reach pinned source;
• the paper’s varlen sort is not attributed to beta23’s main prefix-sum scheduler;
• paper scheduling gains remain source-reported and configuration-specific;
• beta23 is the implementation anchor and beta24 is explicit drift;
• architecture/feature rows are source-admission predicates, not execution proof;
• compile-key fields are grouped by generated-code ownership;
• source, IR, PTX, CUBIN/SASS, loaded symbol, profile, and timing form separate evidence levels;
• FA4 compile time is reconciled without calling it kernel runtime;
• dense, causal, forward, and backward FLOP numerators are derived;
• useful, scheduled, and dense-equivalent throughput are distinguished;
• the main-text B200 and appendix B100 manifests remain unresolved rather than silently merged;
• no paper benchmark is attributed to beta23 or beta24;
• no local GPU artifact, profile, or timing is claimed.

0x0E Failure Lab

Architecture And Dispatch Failures

Failure	Symptom	Root cause	Check
Compile Hopper path for Blackwell and assume compatibility	unsupported instruction or fallback	WGMMA/TMEM contracts differ	record architecture and loaded kernel symbol
API says FlashAttention but generic backend runs	expected speedup absent	shape/dtype/mask unsupported	framework diagnostics plus profiler trace
Compare different causal or deterministic modes	misleading benchmark	semantics and work differ	include complete dispatch key
Treat beta23 source admission as executed support	paper/source table looks complete but deployment fails	validation, compilation, symbol selection, and launch were collapsed	retain artifact and profile selected symbol
Reuse an executor after a missing compile-key field changes	shape-compatible call produces wrong path or stale ABI	generated-code identity is incomplete	mutate every constexpr branch and assert a cache miss
Import beta24 behavior into the beta23 matrix	documentation no longer matches pinned blobs	release drift was treated as retroactive evidence	pin release, commit, file blobs, and re-extract

Pipeline And Synchronization Failures

Failure	Symptom	Root cause	Check
Reuse SMEM stage too early	intermittent wrong output	consumer still reads old stage	barrier phase assertions and stress tests
Reset K/V phase state at a persistent-work boundary	stale wait or deadlock after several tiles	ring generation continuity was discarded	multi-work-tile stage/phase trace
Advance state without releasing K or V	producer hangs on wrap	coordinate motion was confused with ownership transfer	acquire/commit/wait/release conservation counters
Producer and consumer enumerate different sparse blocks	permanent full wait	a payload expected by the consumer was never committed	compare selected block-coordinate manifests
Treat producer tail as a dummy payload	extra consumer wait or false full generation	synchronization closure was confused with data production	assert no payload count changes in tail
Omit the scheduler bootstrap arrival	all consumer groups stop before first scheduled tensor bundle	no named-barrier half-credit starts the ring	source trace plus first-generation barrier oracle
Use the wrong named-barrier participant count	premature release or permanent wait	the two 128-thread consumer halves do not sum to the declared 256	inspect emitted barrier IDs/counts and stress repeated work tiles
Diverge scheduler calls across consumer groups	shape-dependent deadlock	named-barrier generations no longer match	per-group static call ledger for dense/sparse/empty branches
Omit the non-overlap drain arrival	peer or next work tile hangs	final PV consumed the token without passing it	verify J+1 sync/arrive balance per non-empty tile
Re-prime scheduler at every persistent work tile	premature barrier generation or later deadlock	bootstrap token was duplicated instead of preserved	multi-tile token-conservation trace
Wait too early after async issue	correct but slow	overlap collapses into serialization	instruction timeline and stall reasons
Wait too late	race or stale data	dependency not satisfied	sanitizer/reference comparison
Producer gets excessive registers	occupancy or spills	setmaxnreg budget wrong	register count, spills, resident CTAs
Add a third stage blindly	slower kernel	live score tiles exceed register budget	stage-count sweep and occupancy profile
2-CTA pair diverges in mode/order	deadlock or undefined result	cooperative instruction contract violated	paired-CTA control-flow audit

Memory And Layout Failures

Failure	Symptom	Root cause	Check
FP8 V uses FP16 layout assumption	wrong values or unsupported MMA	k-major operand constraint ignored	layout/unit test before MMA
TMEM regions overlap while live	silent corruption	lifetime allocator reused columns early	draw TMEM interval/lifetime chart
DSMEM exchange partitions reduction axis incorrectly	wrong dQ	each CTA lacks full reduction	compare 1-CTA and 2-CTA reference
Larger tile exceeds SMEM/TMEM budget	launch failure or low occupancy	tile choice ignores capacity	static resource ledger plus profiler

Numerical Failures

Failure	Symptom	Root cause	Check
Orthogonal transform differs between Q and K	logits change before quantization	R_Q R_K^T != I	FP64 identity test
Per-tensor scale used where block scale expected	FP8 error spikes on outliers	one scale covers heterogeneous blocks	blockwise error histogram
Polynomial exp used in FP32-sensitive path	accuracy regression	BF16-rounding argument no longer applies	compare before and after target dtype rounding
Conditional rescale threshold too large	overflow or saturated P	reference frame drifts too far	adversarial monotone score rows
Approximate path compared with loose aggregate metric only	rare catastrophic rows hidden	mean error masks tails	max/percentile row and output error

Measurement Failures

report theoretical TFLOP/s without measured time
compare kernels with different tensor shapes
omit compile/warmup overhead inconsistently
include allocation in one path but not another
benchmark a fallback while naming the requested backend
turn a paper's best case into a deployment-wide multiplier
infer semantic long-context quality from kernel throughput
use dense FLOPs for a causal or sparse timing without labeling the convention
call the paper's one-kernel compile speedup a device-runtime speedup
silently normalize the FA4 appendix B100 label into the main-text B200 label
attribute the paper's B200 bars to beta23 or beta24 without a release timing run

0x0F Profiler And Evidence Checklist

For one selected kernel, record:

identity:
    architecture, symbol, source commit, compiler, library version

shape:
    B, S_q, S_k, H_q, H_kv, d_h, mask, causal mode, dtype

capacity:
    registers/thread, SMEM/CTA, TMEM allocation, stages, CTAs/SM

movement:
    HBM bytes, L2 behavior, TMA transactions, SMEM traffic

execution:
    tensor-pipe activity, MUFU activity, issue stalls, barrier stalls

correctness:
    reference output, max/mean error, deterministic-mode behavior

timing:
    warmup, repetitions, distribution, forward/backward boundary

Evidence Hierarchy

weak:
    source code requested a FlashAttention backend

better:
    dispatch diagnostics report eligibility and selected variant

strong:
    profiler identifies the kernel, resource usage, and bottleneck

strongest for a release:
    versioned benchmark + profiler + numerical comparison + fallback test

0x10 What The M4 Execution Branch Now Establishes

M4’s execution/kernel branch closes with a sequence of increasingly concrete resource questions:

M4.1:
    count model parameters, activations, FLOPs, bytes, and phase-specific state

M4.2:
    map tensors onto GPU execution agents, memories, and profiler evidence

M4.3:
    remove quadratic attention intermediates from HBM with tiled online softmax

M4.4:
    improve query-block and warp-level work partitioning

M4.5:
    recover decode parallelism and verify mask/cache/backend contracts

M4.6:
    schedule attention around generation-specific asynchronous engines,
    on-chip memories, low-precision layouts, and moving bottlenecks

M4 also contains a separate empirical branch from M4.1:

M4.7:
    empirical training scaling laws and compute-optimal allocation

M4.8:
    extrapolation, alternative functional forms, and data constraints

M4.9:
    overtraining, downstream metrics, and lifecycle allocation

M4.6 closes the kernel-execution argument; it does not claim that hardware accounting closes empirical model-scaling or lifecycle decisions.

The shared systems rule is:

A mathematical operation does not determine one optimal kernel. The winning schedule depends on tensor shape, numerical mode, persistent state, memory hierarchy, available engines, and the exact hardware generation.

M5.1 starts from the state that survives these kernels across requests: K/V pages, request queues, admission decisions, placement, and scheduler pressure.

The model-wide quantization handoff has five owners: M8.2 chooses calibration, sensitivity, and precision-allocation evidence; M8.3 designs equivalent transforms; M8.6 owns live activation and K/V quantization and scale/state lifetimes; M8.9 validates quantized artifact compatibility and runtime admission; M8.10 owns deployment evaluation, promotion, release, and rollback of the exact candidate. M4.6 retains the attention-local, generation-specific kernel schedule and the evidence it derives; a local FP8 kernel result is not model-wide calibration or release approval.

0x11 Publish Gate

Gate	Status
Hopper resource ceilings are separated from compiled residency	0x02
TMA transaction bytes, full/empty barriers, stage parity, and SMEM liveness are derived	0x02
WGMMA group, proxy-fence, named-barrier, and register-pool domains are distinct	0x02
Pinned CuTe source is traced to expected PTX families without claiming compiled proof	0x02
Paper Algorithm 1 is separated from beta23 Q/K/V pipeline objects and traversal order	0x03
Producer/consumer index, phase, parity, ownership, and conservation are derived	0x03
Non-overlap and peeled intra-warpgroup source paths have legal K/V release points	0x03
Ragged, drain, producer-exit, and persistent-work tails are distinct	0x03
Circular-buffer liveness is closed by a decreasing wait-for proof and failure matrix	0x03
Pingpong benefit is derived from operation/rate imbalance	0x04
Pingpong preserves per-warpgroup numerical state and passes only issue-order permission	0x04
Physical/local warpgroup coordinates, named-barrier IDs, and the 256-thread count are reconciled	0x04
Bootstrap, non-overlap drain, intra-overlap steady state, and persistent token state are distinct	0x04
Cross-warpgroup liveness is closed by token conservation and a decreasing-rank proof	0x04
Paper Figure 1, historical Hopper source, current beta source, and compiled proof remain separate	0x04
Two-stage startup, steady state, drain, and finite timing bound are derived	0x05
QK/PV commit order gives wait_group(1) and wait_group(0) legal meanings	0x05
RS/SS probability placement, register live intervals, and payload limits are separated	0x05
One-, two-, and three-stage variants include compiler and capacity tradeoffs	0x05
Source, PTX, SASS, and dynamic-profile proof claims remain distinct	0x05
FP8 layout and numerical transforms are separated	0x06
Blackwell TMEM and 2-CTA changes are source-grounded	0x07
FA4 forward cycle ledger is derived	0x08
Exponential approximation and rescaling have numerical boundaries	0x09
FA4 backward cycle ledger and DSMEM trade are derived	0x0A
Causal cost, LPT, cache-aware swizzles, static persistence, CLC, and deterministic SPT are separated	0x0B
Paper varlen sorting and beta23 prefix-sum coordinate recovery remain distinct	0x0B
beta23 source admission, compile key, generated artifact, loaded symbol, profile, and timing are separate evidence levels	0x0C
beta24 is recorded as release drift rather than imported into the beta23 support matrix	0x0C
Dense/causal and forward/backward benchmark numerators are derived and timing scopes are separated	0x0D
Main-text B200 and appendix B100 benchmark metadata remain explicitly unresolved	0x0D
Failure diagnosis reaches synchronization, placement, precision, and evidence	0x0E-0x0F
Paper benchmarks remain hardware/version/shape-specific	0x0D
Display equations are balanced and each numerical coordinate has a mechanically checked representative formula	final publication audit
Every local source, article link, figure target, and responsive visual has a checked identity or provenance record	final publication audit
Desktop/mobile light/dark and reduced-motion rendering have no broken image, MathJax error, or article-level horizontal overflow	final publication audit
The local evidence ceiling remains source-audited and unmeasured: no retained GPU artifact, profile, correctness run, or timing sample	0x0C-0x0F

0xFF References

• Local source: Exploring the Transformer Series (19) - FlashAttention v2 and its upgrade
• M4.1 - Transformer Cost Model
• M4.2 - GPU Execution Model for Transformer Kernels
• M4.3 - IO-Aware Attention and FlashAttention-1
• M4.4 - FlashAttention-2 Work Partitioning
• M4.5 - Decode Attention Kernels and Backend Boundaries
• M5.1 - Autoregressive Serving State
• M8.2 - Calibration, Sensitivity, and Precision Maps
• M8.3 - Outlier Geometry and Equivalent Transforms
• M8.6 - Live-Data Quantization: Activations and KV Cache
• M8.9 - Quantized Artifacts and Runtime Integration
• M8.10 - Deployment Evaluation, Release, and Rollback
• Jay Shah et al., FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-precision
• Ted Zadouri et al., FlashAttention-4: Algorithm and Kernel Pipelining Co-Design for Asymmetric Hardware Scaling
• Chee et al., QuIP: 2-Bit Quantization of Large Language Models With Guarantees
• Dao-AILab, FlashAttention fa4-v4.0.0.beta23, commit 2409214a03797b168f648ea30df1adbc09ce658a
• Dao-AILab, FlashAttention fa4-v4.0.0.beta24, commit 849f660f73b176e5ad5670e7f822c7fa9f3eaf8b (release drift audited 2026-07-30; not the implementation anchor)
• NVIDIA, Hopper Tuning Guide
• NVIDIA, Blackwell Tuning Guide
• NVIDIA, Parallel Thread Execution ISA, especially CTA named barriers, cp.async.bulk.tensor, mbarrier, WGMMA, and setmaxnreg
• NVIDIA CUTLASS, cutlass.pipeline API, especially named-barrier participant accounting, pipeline state, acquire/commit, wait/release, and producer tail
• NVIDIA CUTLASS, WGMMA Programming
• NVIDIA CUTLASS, cute.nvgpu.warpgroup API, especially WGMMA register fences, commit groups, and waits
• NVIDIA CUTLASS, tcgen05 MMA Programming
• NVIDIA CUTLASS, CuTe DSL Overview
• NVIDIA CUTLASS, CuTe DSL Code Generation
• NVIDIA CUTLASS, CuTe DSL Debugging And Artifact Retention
• NVIDIA CUTLASS, CuTe DSL JIT Caching
• Quack Kernels, PyPI release history, audited helper versions 0.5.3 and 0.6.1
• Dao-AILab, FlashAttention repository