Skip to content

feat(wan): short-KV text cross-attention Pallas kernel (wan_cross_attn_kernel="pallas") - #509

Open
Perseus14 wants to merge 1 commit into
feat/wan-custom-kernelsfrom
feat/wan-cross-attn-pallas
Open

Perseus14 wants to merge 1 commit into
feat/wan-custom-kernelsfrom
feat/wan-cross-attn-pallas

Conversation

@Perseus14

@Perseus14 Perseus14 commented Oct 15, 2026 •

Copy link
Copy Markdown
Collaborator

Stacked on #488 (feat/wan-custom-kernels). First of four pieces split out of #491. Adds wan_cross_attn_kernel: "pallas" (default "xla" in the Wan YAMLs / wan_runtime_options): 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 (full-row max, exp2, one numerator/denominator, normalised after PV — no online rescaling). It replaces only what would otherwise be the unmasked XLA dot-product fallback. FlaxWanAttention._pallas_cross_attention falls back to XLA (logging the reason once) when the inputs/sharding don't fit, when the mesh is not all v6e/tpu7x (the only generations the 64 MiB VMEM budget was validated on), or when estimate_vmem_bytes exceeds the budget. T2V cross-attention and the text half of I2V cross-attention dispatch to it. In isolation the kernel is 0.51 ms vs 3.94 ms (v6e) and 0.49 ms vs 1.81 ms (tpu7x) per call vs the XLA path inside the model. run_wan_fast_inference.sh turns it on for the v6e/v7 profiles.

E2E latency (TPU v6e-8, Wan 2.2 T2V-A14B, 720p / 81 frames / 40 steps, CP=4, DP=2, warm AOT, DVFS unpinned): 130.6 s generate / 128.4 s denoise (cold AOT: 130.4 s generate / 128.3 s denoise), down from 133.7 s generate / 131.6 s denoise on #488 (−3.1 s generate / −3.2 s denoise end-to-end from Pallas cross-attention alone).

Tests: cross_attention_pallas_test.py (32 test methods / 40 parameterized cases, new; parity grid including the Wan production shard shapes 18900/18944×512 on TPU, guards, VMEM budget logic, FlaxWanAttention dispatch incl. CPU-interpret integration), wan_runtime_options_test.py (+enum validation).

Stack: #477 → #507 → #478 → #508 → #479 → #488 → #509 → #510 → #511 → #491

@Perseus14
Perseus14 requested a review from entrpn as a code owner October 15, 2026 15:54
@github-actions

Copy link
Copy Markdown

@gemini-code-assist gemini-code-assist Bot left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Code Review

This pull request introduces a fused short-KV cross-attention Pallas kernel for TPU, designed to replace the XLA dot-product fallback for short KV sequences in Wan models. It integrates this kernel into FlaxWanAttention, updates configuration files and runtime options to support the new wan_cross_attn_kernel parameter, and adds corresponding unit tests. The review feedback suggests a minor optimization in the Pallas kernel to compute the tile-wide maximum in a single jnp.max call instead of two nested calls.

Comment thread src/maxdiffusion/kernels/cross_attention_pallas.py Outdated
@Perseus14
Perseus14 force-pushed the feat/wan-cross-attn-pallas branch from f78bfda to 4a9cb02 Compare October 15, 2026 17:53
@Perseus14
Perseus14 force-pushed the feat/wan-cross-attn-pallas branch from 4a9cb02 to 1f5539b Compare October 15, 2026 19:20

@mbohlool mbohlool left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The main ask is that a requested pallas raises instead of falling back to XLA, and that CI runs the kernel at least once in interpret mode.

Comment thread src/maxdiffusion/models/attention_flax.py Outdated
Comment thread src/maxdiffusion/tests/cross_attention_pallas_test.py Outdated
Comment thread src/maxdiffusion/configs/base_wan_14b.yml Outdated
Comment thread src/maxdiffusion/kernels/cross_attention_pallas.py Outdated
Comment thread src/maxdiffusion/kernels/cross_attention_pallas.py Outdated
Comment thread src/maxdiffusion/models/attention_flax.py Outdated
@Perseus14
Perseus14 force-pushed the feat/wan-cross-attn-pallas branch from 1f5539b to da317cc Compare October 20, 2026 12:51
@Perseus14
Perseus14 force-pushed the feat/wan-cross-attn-pallas branch 3 times, most recently from cdb5a4d to 56acb00 Compare October 20, 2026 15:13
@Perseus14
Perseus14 force-pushed the feat/wan-cross-attn-pallas branch 2 times, most recently from fc80478 to ed91d0d Compare October 21, 2026 08:48
@Perseus14
Perseus14 force-pushed the feat/wan-cross-attn-pallas branch from ed91d0d to 76b6e85 Compare October 21, 2026 08:55
@Perseus14
Perseus14 force-pushed the feat/wan-cross-attn-pallas branch from 76b6e85 to b7d4ab7 Compare October 21, 2026 09:32
@Perseus14
Perseus14 force-pushed the feat/wan-cross-attn-pallas branch from b7d4ab7 to c086c74 Compare October 22, 2026 08:54
…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`.
@Perseus14
Perseus14 force-pushed the feat/wan-cross-attn-pallas branch from c086c74 to d0a909f Compare October 22, 2026 22:30
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants