Repository navigation
perf(attention): shard-major Ulysses output all-to-all (wan_ulysses_out_a2a="chunked"/"shard_major") - #511
perf(attention): shard-major Ulysses output all-to-all (wan_ulysses_out_a2a="chunked"/"shard_major")#511Perseus14 wants to merge 3 commits into
Conversation
There was a problem hiding this comment.
Code Review
This pull request introduces the wan_ulysses_out_a2a configuration option (supporting "flat", "chunked", and "shard_major" modes) to optimize the inverse Ulysses exchange and avoid expensive post-kernel relayout copies on TPUs. This is achieved by updating the custom splash and ring attention kernels to support 4D shard-major query inputs and output sharding, alongside adding a new test suite to verify correctness. The review feedback highlights a critical bug in _ulysses_ring_custom_attention where use_shard_major fails to check for transpose_out, potentially causing layout mismatches. Additionally, the reviewer suggests using the more idiomatic jnp.square instead of ** 2 in JAX operations and wrapping a long line in ring_attention_kernel.py to comply with PEP 8 line length guidelines.
2cabbb8 to
f12639a
Compare
251437a to
64a7d8a
Compare
f12639a to
86f71b7
Compare
24bbfdf to
f26b420
Compare
86f71b7 to
df7da19
Compare
mbohlool
left a comment
There was a problem hiding this comment.
The main asks: a requested shard_major should not silently run flat, and CI should run the single-device kernel test.
df7da19 to
808d7bd
Compare
388384b to
97c7129
Compare
808d7bd to
01c2d43
Compare
97c7129 to
cf59fab
Compare
6f8ec7d to
5b4eddb
Compare
cf59fab to
2bd2d3b
Compare
2bd2d3b to
72712b6
Compare
5b4eddb to
52ec396
Compare
72712b6 to
c99f1c3
Compare
208901a to
09d9dab
Compare
c99f1c3 to
1924230
Compare
09d9dab to
84a68e3
Compare
1924230 to
74f32d8
Compare
84a68e3 to
d98f3ab
Compare
74f32d8 to
59c1ee5
Compare
d98f3ab to
da01a58
Compare
59c1ee5 to
8697341
Compare
8697341 to
b31bd66
Compare
da01a58 to
2ea7bb9
Compare
…n_kernel="pallas")
First of four pieces split out of the Wan fast-inference optimizations PR.
The switch defaults off ("xla") in the Wan YAMLs and in wan_runtime_options;
it is a plain config key plumbed to FlaxWanAttention through
`attention_config_entries(config)`.
- `kernels/cross_attention_pallas.py` keeps a head block's whole 512-token text
KV resident in VMEM and computes an exact single-pass softmax per query tile
(per-row max, exp2, MXU sum reduction, normalised after PV; no
running max / online rescaling). It replaces only what would otherwise be
the unmasked XLA dot-product fallback (`_would_use_unmasked_dot_product`,
which shares `_FLASH_MIN_SEQ_LENGTH_KERNELS` with the dispatcher).
- `wan_cross_attn_kernel="pallas"` fails closed (`ValueError`): static
preconditions (TPU mesh or `wan_cross_attn_cpu_interpret=True`,
`dim_head % NUM_LANES == 0`, compatible `attention_kernel`,
`not use_memory_efficient_attention`, and validated TPU VMEM generation) are
checked in `FlaxWanAttention.__init__`, and dynamic input/sharding/VMEM
preconditions are checked in `FlaxWanAttention._pallas_cross_attention`.
- Dispatch: T2V cross-attention, and the text half of I2V cross-attention
(both the image-token split and the text-only path), call
`_pallas_cross_attention` first (which returns `None` only when
`wan_cross_attn_kernel == "xla"`). The sharded wrapper is cached
(`_build_sharded_pallas_cross_attention`) and wrapped with
`with_xla_backward` since `pallas_call` has no transpose rule.
- `wan_cross_attn_cpu_interpret` (default False, tests only) lets the kernel
run under Pallas interpret mode in CI (both on CPU and on unvalidated TPU
runners).
- In isolation the kernel is 0.51 ms vs 3.94 ms (v6e) and 0.49 ms vs 1.81 ms
(tpu7x) per call for the XLA path inside the model.
- `run_wan_fast_inference.sh`: the v6e and v7 profiles set
`WAN_CROSS_ATTN_KERNEL=pallas` (overridable); the generic profile keeps `xla`.
- `wan_block_benchmark.py` passes `split_head_dim` (defaulting to `False` to
match the Wan YAMLs), `ulysses_attention_chunks`, and `use_k_centering` from
the config to the benchmarked block.
Tests
- `cross_attention_pallas_test.py` (31 test methods, 32 parameterized cases):
golden parity (ragged tiles, `k_prescaled`, batch), fail-closed guards, VMEM
budget resolution / working-set estimate / unusable reasons, and
`FlaxWanAttention` dispatch (interpret-mode integration, precondition
`ValueError` checks, I2V text-vs-image split).
- `wan_runtime_options_test.py` (12): YAML default and enum validation.
- CI: 30 of the 32 cases in `cross_attention_pallas_test.py` (all parity,
guard, VMEM budget, interpret-mode integration, and I2V split tests) run in
GitHub Actions; only the full-shard compiled TPU tests
(`test_wan_production_shard_shapes` and
`test_pallas_mode_uses_the_kernel_and_matches_xla`) skip in CI.
- `end_to_end/tpu/run_wan_stack_tests.sh`: adds `cross_attention_pallas_test.py`.
Second of four pieces split out of the Wan fast-inference optimizations PR.
The switch defaults off ("conv") in the Wan YAMLs and in wan_runtime_options;
it is a plain config key plumbed to WanModel through
`attention_config_entries(config)` and resolved at build time
(`_wan_opt`) so it lives on the GraphDef.
- With `"tokens"`, the Conv3d patch embedding becomes a patchify + matmul on
the token-sharded layout (`embed_patches_as_tokens`): the non-overlapping
patches (stride == kernel) are flattened in the kernel's
`(p_t, p_h, p_w, C)` order, sharded by token and multiplied by the
replicated `[p_t*p_h*p_w*C, D]` kernel (with kernel and bias cast to
`hidden_states.dtype` to match `nnx.Conv.promote_dtype`). Same result as
`collapse(conv(x), 1, -1)` up to accumulation order, but produced directly
in the layout the transformer blocks use, so XLA no longer emits the
activation all-gather / collective-permute when collapsing `(F, H, W)`
into the context-sharded token axis (12.8 ms per step on tpu7x, exposed).
- `can_embed_patches_as_tokens` checks the conv really is a plain
stride-equals-kernel 3D patchify (no groups, masks, dilation, and `'SAME'`
pads nothing). When `wan_patch_embed_mode="tokens"` is requested on an
indivisible latent grid, `WanModel.__call__` raises `ValueError` instead of
silently falling back to `nnx.Conv`.
- `run_wan_fast_inference.sh`: every profile sets `WAN_PATCH_EMBED_MODE=tokens`
(overridable).
Tests
- `wan_patch_embed_test.py` (8): tokens vs conv parity (function level and
through `WanModel`), the eligibility predicate, fail-closed `ValueError` on
indivisible latent grids, and the config-driven mode. All 8 tests run
unconditionally in GitHub Actions.
- `wan_runtime_options_test.py` (12): the new switch's YAML default and enum
validation; `wan_patch_embed_mode` is excluded from the `FlaxWanAttention`
attribute check since `WanModel` consumes it.
- `end_to_end/tpu/run_wan_stack_tests.sh`: adds `wan_patch_embed_test.py`.
…ut_a2a="chunked"/"shard_major")
Third of four pieces split out of the Wan fast-inference optimizations PR.
The switch defaults off ("flat") in the Wan YAMLs and in wan_runtime_options;
it is a plain config key (`"flat"` | `"chunked"` | `"shard_major"`) plumbed to
FlaxWanAttention / NNXAttentionOp / AttentionOp and down to the Ulysses
kernels through the dispatch context.
- `_ulysses_seq_to_heads` is the inverse Ulysses exchange (attention output,
head-sharded -> sequence-sharded), and validates `seq_len % num_shards == 0`
up front for all modes. `"flat"` is the previous
`all_to_all(split_axis=seq, concat_axis=1)`; XLA reorders the whole output
so the sequence axis is major before the collective and back again after,
four relayout copies per layer at the 720p shard shape (~1.5 ms on v6e,
~0.7 ms on tpu7x). `"chunked"` / `"shard_major"` uses a single
layout-preserving `all_to_all` over a reshaped `(num_shards, chunk)` leading
axis, so no relayout copies are inserted (also wired into the non-custom
`attention="ulysses"` branch). Same data movement, same values.
- With the custom splash kernel (`S_local % block_q == 0`, `D == 128`,
`heads_per_tile == 1`), the forward Q exchange also stays shard-major (4-D q,
`[B, U, H/U, S_local, D]`) and the kernel writes the shard-major output
directly (`out_num_shards`), on both the 1D Ulysses path (v6e) and the 2D
Ulysses+Ring path (tpu7x): `custom_splash_attention.py` gains a 4-D q
BlockSpec and a 4-D shard-major output for `_splash_attention_forward` and
`_splash_attention_forward_ring`; `ring_attention_kernel.py` threads
`out_num_shards` through `_custom_ring_attention_forward` /
`make_custom_ring_attention` and sizes the o/l/m accumulators accordingly
(`save_residuals` with `out_num_shards > 1` raises `NotImplementedError`:
the lse contract stays `(num_q_heads, orig_q_seq_len)`).
`_compute_fixed_m_metadata` accepts the 5-D shard-major query.
- `"shard_major"` vs `"chunked"`: `"shard_major"` is the strict mode that
requires the 4-D `q` / `out_num_shards` Pallas kernel path and raises
`ValueError` when any gate precondition (`use_custom_kernel`,
`S_local % block_q == 0`, `D == 128`, `heads_per_tile == 1`,
`not transpose_out`, `attention_mask is None`) fails; `"chunked"` is the
lenient mode that takes the 4-D kernel path when eligible and falls back to
the chunked inverse `all_to_all` in `_ulysses_seq_to_heads` otherwise.
- Reachability: at 720p/81f with CP=4 the unpadded per-device sequence is
18,900 tokens, which no launcher `block_q` divides (v6e `BQ=9472`, v7
`BQ=6400` at this PR), so at this PR the launcher's `v6e`/`v7` profiles set
`DEFAULT_ULYSSES_OUT_A2A=chunked`; once `feat/wan-fast-inference-optimizations`
pads the sequence to 18,944 (`wan_seq_pad=lane`) and sets v7 `BQ=4736`
(`9472 * 2 = 18,944` on v6e, `4736 * 4 = 18,944` on v7), it switches both
profiles to `DEFAULT_ULYSSES_OUT_A2A=shard_major`.
Tests
- `ulysses_out_a2a_test.py` (8): flat vs chunked vs shard_major parity of the
inverse exchange, the mode guard, the `seq_len % num_shards != 0` guard, the
`save_residuals + shard-major` guard, the strict `"shard_major"` kernel-gate
`ValueError` guard, the 4-D q ring kernel bit-exact vs 3-D (running in CI via
`interpret=True`), and bit-exact wrapper output for the shard-major 1D Ulysses
and 2D Ulysses+Ring paths vs the flat layout (spy-asserting
`out_num_shards == U`).
- `wan_runtime_options_test.py` (12): the new switch's YAML default and enum
validation.
- `end_to_end/tpu/run_wan_stack_tests.sh`: adds `ulysses_out_a2a_test.py`.
2ea7bb9 to
8f26fae
Compare
b31bd66 to
31ec68f
Compare
Stacked on
feat/wan-token-patch-embed. Third of four pieces split out of #491. Addswan_ulysses_out_a2a: "chunked" | "shard_major"(default"flat"):_ulysses_seq_to_headsperforms the inverse Ulysses exchange as a single layout-preservingall_to_allover a reshaped(num_shards, chunk)leading axis, so XLA stops inserting the four relayout copies per layer that the flatsplit_axis=seqform costs at the 720p shard shape (~1.5 ms on v6e, ~0.7 ms on tpu7x). Same data movement, same values. With the custom splash kernel (S_local % block_q == 0,D == 128,heads_per_tile == 1) the forward Q exchange also stays shard-major (4-Dq) and the kernel writes the shard-major output directly (out_num_shards), on both the 1D Ulysses path (v6e) and the 2D Ulysses+Ring path (tpu7x):custom_splash_attention.pygains a 4-DqBlockSpec/ 4-D output,ring_attention_kernel.pythreadsout_num_shardsthrough the ring accumulators (save_residuals=Truewithout_num_shards > 1raisesNotImplementedError), and_compute_fixed_m_metadataaccepts the 5-D shard-major query.Reachability note: The Pallas shard-major Q/out path requires
S_local % block_q == 0. At 720p/81f withCP=4the per-device sequence is 18,900 tokens, which no launcherblock_qdivides (BQ=9472on v6e,BQ=6400on v7 at this PR), so on this PR alone the launcher'schunkedsetting only changes the inverse all-to-all; the 4-Dq/out_num_shardsPallas kernel path engages once#491(feat/wan-fast-inference-optimizations) pads the sequence to 18,944 (wan_seq_pad=lane) and setsBQ=4736(4736 × 4 = 18,944).E2E latency (TPU v6e-8, Wan 2.2 T2V-A14B, 720p / 81 frames / 40 steps,
CP=4, DP=2,wan_ulysses_out_a2a="chunked", DVFS unpinned): 128.9 s generate / 126.7 s denoise (29.6 s compile), down from 130.6 s generate / 128.5 s denoise on #510 (−1.7 s generate / −1.8 s denoise end-to-end from the layout-preserving inverse Ulysses all-to-all alone).Tests:
ulysses_out_a2a_test.py(6, new; flat vs chunked vs shard_major parity, mode guard,save_residuals+ shard-major guard, bit-exact 4-D ring kernel output, and bit-exact wrapper output on the 1D and 2D paths asserting via spy thatout_num_shards == U),wan_runtime_options_test.py(+enum validation).Stack: #477 → #507 → #478 → #508 → #479 → #488 → #509 → #510 → #511 → #491