ci(build): bump FA4, quack and ffpa-attn - #3733
Merged
Merged
Conversation
The pinned flash-attention ref 002cce0a1 (2026-07-03) predates both CuTe compile-key fixes, so every forward/backward with a tensor max_seqlen -- i.e. every varlen/packed-sequence step -- rebuilt the compile key and recompiled the kernel. Reported against a Qwen3.5-VL packed-sequence finetune on B300. Bump to ce088ab9ce0f (main), which adds over the old pin: - #2507 [CuTe, Bwd] fix backward compile key churn (max_seqlen tensor) - #2762 [CuTe, Fwd] stabilize tensor max_seqlen compile key - #2745 [CuTe] fix forward dynamic-shape correctness - #2819 [CuTe] speed up scalar SM100 mask compilation #2798 raised the cute requirement to nvidia-cutlass-dsl>=4.6.2, so pin the FA4 CUTLASS DSL install to 4.6.2 to match. Kept exact rather than floating so the image stays reproducible. Note this ref also feeds the FA3 (Hopper) wheel, which is built by default, so the bump affects the default x86 image and not just INSTALL_FA4=true builds. Signed-off-by: Alexandros Koumparoulis <akoumparouli@nvidia.com>
Signed-off-by: Alexandros Koumparoulis <akoumparouli@nvidia.com>
Contributor
Author
|
/ok to test 325d1ed |
Contributor
Author
|
Runtime validation update:
The separate native text-path/config work is stacked in #3734. |
…o 0.6.4 flash_attn.cute (FA4) at the pinned FLASH_ATTN_REF builds against nvidia-cutlass-dsl 4.6.2. quack-kernels and ffpa-attn both import the same `cutlass` package, so all three have to agree on one version. - quack-kernels 0.6.1 -> 0.6.4 (pins cutlass-dsl 4.6.2; 0.6.1 pinned 4.6.0) - ffpa-attn 0.2.2 -> 0.2.3 (first release pinning cutlass-dsl 4.6.2 and quack-kernels 0.6.4; 0.2.2 hard-pinned 4.6.0/0.6.1) - ffpa extra's explicit cutlass-dsl pin 4.6.0 -> 4.6.2 Because ffpa-attn 0.2.3 already declares the versions FA4 needs, the resolution is conflict-free -- no override-dependencies entries required. Corrects the apache-tvm-ffi note as well: FA4 declares apache-tvm-ffi>=0.1.12 but never imports tvm-ffi, so the <=0.1.11 cap that keeps tilelang working is not mutually exclusive with FA4. Both lock files regenerated with uv 0.8.22 (the version CI pins); the diff is confined to the six bumped packages. Signed-off-by: Alexandros Koumparoulis <akoumparouli@nvidia.com>
Contributor
Author
|
/ok to test 9fe6cfd |
akoumpa
force-pushed
the
akoumparouli/bump-fa4-pin
branch
from
August 29, 2026 20:32
924bb0f to
9fe6cfd
Compare
Contributor
Author
|
all tests passed in https://github.com/NVIDIA-NeMo/Automodel/actions/runs/33269762967/job/99153796598?pr=3733 ; waiting for nemoci |
Flips the default in both stages that declare the ARG: wheel_builder (gates building the flash_attn/cute wheel) and automodel_final (gates the nvidia-cutlass-dsl[cu13] install and the flash_attn/cute symlink). Docker does not inherit ARGs across stages, so flipping only the first would build the wheel while skipping the DSL install and the symlink -- a half-installed FA4. Both FA4 wheels are installed --no-deps, so flash-attn-4's declared apache-tvm-ffi>=0.1.12 bound is never resolved and the <=0.1.11 cap that keeps the tilelang kernels working still holds. Signed-off-by: Alexandros Koumparoulis <akoumparouli@nvidia.com>
Contributor
Author
|
/ok to test cd642d4 |
akoumpa
enabled auto-merge (squash)
August 30, 2026 16:41
The note claimed flash-attn-4 never imports tvm-ffi. That is wrong: flash_attn/cute/cache_utils.py has a bare top-level `import tvm_ffi`, reached eagerly through flash_attn.cute.__init__ -> interface.py, so importing flash_attn.cute hard-requires it. The cap itself is still right, but for a different reason. FA4 is installed --no-deps, so its declared >=0.1.12 floor is never resolved, and the tvm_ffi that is present comes via quack-kernels (>=0.1.6,<0.2) pinned here to 0.1.11. FA4 therefore runs against an ffi older than it declares; the two symbols it references, tvm_ffi.Function and tvm_ffi.__version__, both exist in 0.1.11. Note this is unexercised: nothing selects attn_implementation="flash_attention_4" until the FA4 backend lands, so flash_attn.cute is never imported by any test here. Worth knowing given docker/Dockerfile sets FLASH_ATTENTION_CUTE_DSL_CACHE_ENABLED=1, which turns on the AOT cache path in that same cache_utils.py. Verified against flash-attention ce088ab9, the pinned FLASH_ATTN_REF. Comment only; no dependency or lock change. Signed-off-by: Alexandros Koumparoulis <akoumparouli@nvidia.com>
Contributor
Author
|
(i'll fm since the last commit is just a comment change). |
akoumpa
disabled auto-merge
August 30, 2026 20:16
akoumpa
added a commit
that referenced
this pull request
Sep 1, 2026
Resolve conflicts with #3733, which landed the same FA4 enablement independently: - docker/Dockerfile: both sides set INSTALL_FA4=true, FLASH_ATTN_REF=ce088ab9 and cutlass-dsl 4.6.2. Kept this branch's rationale comments, corrected to match the verified behavior -- flash_attn/cute/cache_utils.py does import tvm_ffi at the pinned ref, it just only touches symbols apache-tvm-ffi 0.1.11 already provides. - pyproject.toml: took main's apache-tvm-ffi comment (accurate on the same point) and its ffpa entry. ffpa-attn 0.2.3 pins cutlass-dsl 4.6.2 / quack-kernels 0.6.4 natively, so the override-dependencies workaround this branch added for 0.2.2 is no longer needed; dropped it here and from docker/common/uv-pytorch.toml. - uv.lock, docker/common/uv-pytorch.lock: main's, matching the resolved pyproject. Kept the update_pyproject_pytorch.sh sed-range fix (anchor the delete range at a column-0 ']'), reworded since the entry it cited is gone.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
What
Bumps the pinned
flash-attentionref used to build the FA3 (Hopper) and FA4 (Blackwell CuTe) wheels:and the FA4 CUTLASS DSL install
nvidia-cutlass-dsl[cu13]==4.6.0.dev0->==4.6.2.It also enables FA4 by default (
INSTALL_FA4false->true) and moves the Python dependencyset onto that same CUTLASS DSL, so everything in the image agrees on one
cutlass:nvidia-cutlass-dsl4.6.0.dev0/4.6.04.6.2docker/Dockerfile,pyproject.toml(ffpaextra)quack-kernels0.6.10.6.4pyproject.toml(core dependency)ffpa-attn>=0.2.2>=0.2.3pyproject.toml(ffpaextra)uv.lockanddocker/common/uv-pytorch.lockwere regenerated with uv 0.8.22 (the version CI pins).The lock diff is confined to those three packages plus the four
nvidia-cutlass-dsl-libs-*splitsthey pull -- seven in total, no marker churn.
Why
The old pin
002cce0a1dates from 2026-07-03 and predates the CuTe compile-key fixes. On that ref, any call with a tensormax_seqlen— i.e. every varlen / packed-sequence step — rebuilds the compile key, so the kernel recompiles on every call instead of hitting the compile cache.This was reported externally against a Qwen3.5-VL packed-sequence full finetune on B300 (
examples/vlm_finetune/qwen3_5/qwen3_5_4b_neat_packing.yaml,pack_size16384), where it shows up as per-step recompilation stalls.Commits gained over the old pin:
[CuTe, Bwd]Fix backward compile key churn due to pickling, max_seqlen is a tensor[CuTe, Fwd]Stabilize tensor max_seqlen compile key[CuTe]Fix forward dynamic-shape correctness[CuTe]Speed up scalar SM100 mask compilation29 upstream commits / 67 files total in the range.
CUTLASS DSL pin
#2798 raised
flash_attn/cute's requirement fromnvidia-cutlass-dsl==4.6.0.dev0to>=4.6.2, so the Dockerfile's DSL install has to move with the ref or FA4 runs against a DSL older than it declares. Pinned exactly at4.6.2(the declared floor, and[cu13]is published for it) rather than left floating, to keep the image reproducible.The Dockerfile alone was not enough, because
quack-kernelsandffpa-attnimport the samecutlasspackage and pinned it independently.quack-kernels0.6.1 pinned4.6.0;ffpa-attn0.2.2 hard-pinned
4.6.0andquack-kernels==0.6.1. Both had to move for the resolution toadmit
4.6.2.ffpa-attn0.2.3 is the first release that pinsnvidia-cutlass-dsl==4.6.2andquack-kernels==0.6.4-- exactly what FA4 needs -- so the conflict resolves on its own and nooverride-dependenciesentries are required. Do not relax the>=0.2.3floor.Note that
ffpa-attnhas no runtime surface in CI: the image builds withAUTOMODEL_INSTALL=all,and the
allextra does not includeffpa. That bump only affects resolution.Blast radius
FLASH_ATTN_REFfeeds both the FA3 and FA4 builds, andINSTALL_FA3=trueis the default on x86 — so this changes the default x86 image, not onlyINSTALL_FA4=truebuilds. That is the main thing CI needs to confirm. The range touches Hopper (#2746removes SM100 functions from hopper,#2756block-sparse bwd wait on SM90), which is why this wants a full pipeline rather than a smoke test.This PR does change which image gets FA4:
INSTALL_FA4flipsfalse->true. The ARG isdeclared in two stages and Docker does not inherit ARGs across stages, so both were flipped --
wheel_builder(gates building theflash_attn/cutewheel) andautomodel_final(gates thenvidia-cutlass-dsl[cu13]install and theflash_attn/cutesymlink). Flipping only the firstwould leave FA4 half-installed: the wheel present, the DSL and symlink missing.
That means the default image now ships FA4 on both arches. arm64 skips FA3 (
TARGETARCHguard) but not FA4, so the arm build now runs the
nvidia-cutlass-dsl[cu13]==4.6.2install --manylinux_2_28_aarch64wheels are published for it, for bothlibs-cu13andlibs-base.Both FA4 wheels install with
--no-deps(pip wheel --no-deps, thenpip install --no-deps), soflash-attn-4's declared
apache-tvm-ffi>=0.1.12,<0.2is never resolved and the<=0.1.11cap thatkeeps the tilelang kernels working still holds, so FA4 and tilelang can share one image.
FA4 does import
tvm_ffi(flash_attn/cute/cache_utils.py, eagerly viaflash_attn.cute.__init__),so it runs against an ffi older than it declares. The two symbols it references,
tvm_ffi.Functionandtvm_ffi.__version__, both exist in 0.1.11. This is unexercised here --nothing selects
flash_attention_4yet -- and worth knowing because the Dockerfile setsFLASH_ATTENTION_CUTE_DSL_CACHE_ENABLED=1, which enables the AOT cache path in that same file.Fallback
If the range proves too wide, the conservative stopping point is
69e1bcbe7(2026-08-07) — keeps #2507 + #2762 + #2745 and needs no CUTLASS DSL change, giving up only #2819.Testing
All runs are against
cd642d497(FA4 enabled). Each pipeline builds a fresh container from thisbranch, so the regenerated
uv.lockis exercised end to end. Each eos scope was run twice on thesame commit, so a recipe that fails in one and passes in the other is noise rather than a regression.
performance, folderbenchmarknightly, all foldersperformance, folderbenchmarkqwen2_5_7b_quack+_baselinepass on SM100maincontrol on the real failuresGitHub Actions "CICD NeMo": success on
cd642d497.Container builds are green on both arches:
uv sync --lockedinstalls the regenerated lock, theFA3 source build against the new
FLASH_ATTN_REFcompiles (x86), and the FA4 wheel plusnvidia-cutlass-dsl[cu13]==4.6.2install on x86 and arm.Why Blackwell was required
FA4's CuTe kernels are SM100-targeted and are never reached on H100, and cutlass-dsl / quack dispatch
to different codegen per arch (
gemm_sm90vsgemm_sm100). An H100-only run would leave the point ofenabling FA4 untested. The arm image is the sharper test: arm64 skips FA3 entirely, so FA4 is the only
flash backend present there.
Failure accounting
No failure in any FA4-on run is attributable to this PR. Every one resolves to:
main—nemotron_nano_9b_squad,nemotron_nano_9b_squad_peft,nemotron_super_v3_hellaswag_peftall fail onmaintoo (control 65263128), withCHECKPOINT_ROBUSTNESS_PHASE_FAILURE.gpt_oss_20bandnemotron_flash_1b_squadeach fail once and pass onmainplus theduplicate run.
curl: (5) Could not resolve proxy, ~65 swalltime) and one runner-workspace fault (
getcwd: cannot access parent directories). Both alsooccur on the FA4-off runs.
*_vllm_deployjobs fail on a checkpoint their upstream test never wrote.The three GB200 failures are environmental, not kernel:
qwen3_moe_30b_quack_deepephitsnon_pp_size=4 must be a multiple of ep_size=8(the recipe hardcodesep_size: 8; GB200 nodes have4 GPUs), and both
llama3_1_8b_quack_ropevariants fail identically on a cold HF cache(
Couldn't instantiate the backend tokenizer) — the quack variant and its baseline failing the sameway rules out anything quack-specific.
automodel_{llm,vlm}_benchmark_testsreporting failed on the nightly-scope runs is the knownempty-child case (the benchmark folders ship no
nightly_recipes.yml); it reproduces onmain.Known gap
No test on this branch executes an FA4 kernel. Nothing selects
attn_implementation="flash_attention_4"until the FA4 backend lands in #3734. The
INSTALL_FA4=trueflip is therefore validated as builds,installs and breaks nothing — on both arches — rather than as FA4 produces correct results. The
latter is covered by #3734's
tests/functional_tests/attention/test_fa4_packed_parity.py.nemo-ci fixes required to reach GB200
Two nemo-ci bugs blocked Blackwell entirely and are fixed in separate MRs:
oci_hsgnever setGRES, so sbatch rejected every submission(
Cannot find GPU specification).NPROC_PER_NODEcame from a hardcoded cluster list missingoci_hsgandhecate,so torchrun spawned 8 ranks onto 4-GPU nodes (
CUDA error: invalid device ordinal).