Skip to content

feat(speculative): add Domino online training path on top of DFlash - #2819

Merged
HuiyingLi merged 3 commits into
NVIDIA-NeMo:mainfrom
khazic:khazic/feat/dflash-domino-training
Jun 28, 2026
Merged

feat(speculative): add Domino online training path on top of DFlash#2819
HuiyingLi merged 3 commits into
NVIDIA-NeMo:mainfrom
khazic:khazic/feat/dflash-domino-training

Conversation

@khazic

@khazic khazic commented Jun 28, 2026

Copy link
Copy Markdown
Contributor

Objective

Add a Domino online training path on top of the existing DFlash speculative-decoding stack, ported from SpecForge (sgl-project/SpecForge#571). The existing DFlash path is left unchanged.

Domino extends the parallel DFlash draft backbone with a lightweight causal correction head. DFlash drafts a whole block in a single non-causal forward, so each predicted position is blind to the (drafted) tokens earlier in its own block. The Domino head fixes that: a GRU encodes a causal state from the block's previous tokens, and a low-rank projection of [backbone hidden | GRU state] produces a logit delta that is added to the parallel base logits.

Design

Training jointly supervises two logits with a base-anchor curriculum:

loss = (1 - lambda_base) * final_loss + lambda_base * base_loss

final_loss is the CE on the Domino-refined logits, base_loss the CE on the backbone-only base logits. lambda_base decays linearly to 0 over the first lambda_base_decay_ratio fraction of training, so early steps keep the parallel backbone strong and later steps let the correction head take over.

The integration reuses DFlash end to end:

  • DominoTrainerModule subclasses DFlashTrainerModule, reusing anchor sampling, the noise-block construction, and the block attention mask. Only the head, the dual-logit loss, and the metrics are Domino-specific.
  • TrainDominoRecipe subclasses TrainDFlashRecipe through four behavior-preserving seam methods added to the DFlash recipe (_build_dflash_config, _build_trainer_module, _run_trainer_step, _log_extra_train_metrics). The DFlash recipe behavior is unchanged.
  • The Domino head (prefix_gru, embed_proj) lives on the DFlash draft model, gated by dflash_config.projector_type='domino', so a pure-DFlash draft has no extra parameters.

RoPE precision fix

While validating, I found that model.to(bfloat16) rounds the Qwen3 rotary inv_freq buffer to bf16, whereas HF from_pretrained and the SGLang serving runtime keep it in fp32. That train/inference RoPE mismatch grows with absolute position and erodes draft acceptance. The draft now recomputes fresh fp32 frequencies from the rotary config after any low-precision cast (a bf16 round-trip cannot be undone by upcasting), covering the training build and reload paths. This benefits DFlash and Domino alike.

Files

File Change
nemo_automodel/components/speculative/dflash/domino_core.py New DominoTrainerModule, dual-logit loss, metrics, get_lambda_base
nemo_automodel/components/speculative/dflash/draft_qwen3.py Optional Domino head + fp32 RoPE pin
nemo_automodel/recipes/llm/train_dflash.py Four behavior-preserving seams
nemo_automodel/recipes/llm/train_domino.py New TrainDominoRecipe
examples/speculative/dflash/qwen3_domino.yaml Example config
tests/unit_tests/... test_domino_core.py, test_train_domino.py, RoPE-fp32 regression tests

Validation

Unit tests (full DFlash + Domino + speculative suite): 393 passed, 1 skipped on an A800 (including new RoPE-fp32 regression tests).

End-to-end smoke (Qwen3-0.6B target, ShareGPT-style data, sdpa backend): a full run to the final checkpoint. Loss trends down, final_loss separates below base_loss, lambda_base decays on schedule, and the saved consolidated safetensors carries the Domino head (prefix_gru.*, embed_proj.*) with projector_type: domino in the config.

Usage

automodel examples/speculative/dflash/qwen3_domino.yaml --nproc-per-node 8

The example expects the same target-model and dataset setup as the DFlash online-training examples, with the Domino architecture and curriculum parameters read from the YAML.

References

khazic added 3 commits June 28, 2026 15:51
Domino (sgl-project/SpecForge#571) extends the parallel DFlash draft backbone with a lightweight causal correction head: a GRU encodes a causal state from each block's previous tokens, and a low-rank projection of [backbone hidden | GRU state] adds a correction to the parallel base logits. Training jointly supervises the Domino-refined and backbone-only logits with a base-anchor curriculum (loss = (1-lambda_base)*final_loss + lambda_base*base_loss, lambda_base decaying to 0).

The DFlash path is unchanged: the draft model gains an optional Domino head (projector_type='domino'), a new DominoTrainerModule reuses DFlash anchor sampling and the block mask, and TrainDominoRecipe subclasses the DFlash recipe via four behavior-preserving seams. Adds an example YAML and unit tests (domino core + recipe seams).

Signed-off-by: khazic <khazzz1c@gmail.com>
model.to(bfloat16) rounded the rotary inv_freq buffer to bf16, but HF from_pretrained and the SGLang serving runtime keep it in fp32. That train/inference RoPE mismatch grows with absolute position and erodes draft acceptance. Recompute fresh fp32 frequencies from the rotary config after any low-precision cast (a bf16 round-trip cannot be undone by upcasting), covering the training build and reload paths.

Signed-off-by: khazic <khazzz1c@gmail.com>
…nt mask clone

/simplify cleanups (behavior unchanged): _compute_weighted_losses returns only the three losses instead of also re-returning flattened views (the caller recomputes them as views where needed), and the pre-decay eval mask references weight_mask directly since the decay step rebinds it out-of-place, avoiding a per-step clone.

Signed-off-by: khazic <khazzz1c@gmail.com>
@khazic
khazic requested a review from a team as a code owner June 28, 2026 08:53
@copy-pr-bot

copy-pr-bot Bot commented Jun 28, 2026

Copy link
Copy Markdown

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

@HuiyingLi

Copy link
Copy Markdown
Contributor

/claude review

@claude claude Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

LGTM

@HuiyingLi

Copy link
Copy Markdown
Contributor

/ok to test 0c29737

@HuiyingLi
HuiyingLi merged commit 1bd6828 into NVIDIA-NeMo:main Jun 28, 2026
148 checks passed
akoumpa added a commit that referenced this pull request Aug 30, 2026
* build: bump FA3/FA4 pin to pick up CuTe compile-key fixes

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>

* build: pin full flash-attention commit

Signed-off-by: Alexandros Koumparoulis <akoumparouli@nvidia.com>

* build: bump cutlass-dsl to 4.6.2, ffpa-attn to 0.2.3, quack-kernels to 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>

* build: enable FA4 (INSTALL_FA4=true) by default

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>

* build: correct the apache-tvm-ffi comment

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>

---------

Signed-off-by: Alexandros Koumparoulis <akoumparouli@nvidia.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants