Skip to content

feat(speculative): context parallelism for draft-model training#2918

Merged
HuiyingLi merged 34 commits into
NVIDIA-NeMo:mainfrom
kashif:kashif/feat/draft-cp
Jul 11, 2026
Merged

feat(speculative): context parallelism for draft-model training#2918
HuiyingLi merged 34 commits into
NVIDIA-NeMo:mainfrom
kashif:kashif/feat/draft-cp

Conversation

@kashif

@kashif kashif commented Jul 3, 2026

Copy link
Copy Markdown
Contributor

Adds context parallelism (cp_size>1) to the DFlash/Domino/JetSpec and DSpark draft trainers for long-context training, and fixes the same path in EAGLE-3.

The approach follows EAGLE-3's existing seam: only the frozen target's forward is sharded across the cp mesh, and its captured hidden states are gathered back to the full sequence before the draft runs — so the draft's block/anchor masks are never sequence-split. The draft shards over the dp submesh; cp ranks in a dp group share data and draft weights.

Two things that bit me and are worth a careful look in review:

  • The drafts sample anchor positions from the global RNG, so cp (and tp) replicas would diverge if seeded per global rank. Now seeded per dp rank so they stay in sync.
  • torch's context_parallel ring doesn't actually fire for a plain HF target — q/k/v reach SDPA as local tensors, so each rank silently attended only to its own sequence shard. Since the target is frozen (forward-only), I all-gather K/V and attend the local Q against the full sequence with a global causal mask (new attach_cp_kv_gather_hooks). This needs the HF torch-SDPA path, so cp_size>1 requires target_force_hf=true.

Validated on 2 GPUs (Qwen3-0.6B): with the fix, per-step draft loss matches the 1-GPU baseline within bf16 noise (mean signed +0.18%, symmetric above/below), versus a systematic +2–7% before. Unit tests pass.

Defaults (cp_size=1) are unchanged, and DSpark CP is gated to the dense Qwen3 target (the MoE targets already run under their own EP/FSDP mesh).

kashif added 6 commits July 3, 2026 21:16
Signed-off-by: Kashif Rasul <kashif.rasul@gmail.com>
Signed-off-by: Kashif Rasul <kashif.rasul@gmail.com>
Signed-off-by: Kashif Rasul <kashif.rasul@gmail.com>
…n sync

Signed-off-by: Kashif Rasul <kashif.rasul@gmail.com>
…ence

Signed-off-by: Kashif Rasul <kashif.rasul@gmail.com>
Signed-off-by: Kashif Rasul <kashif.rasul@gmail.com>
@kashif kashif requested a review from a team as a code owner July 3, 2026 23:06
@copy-pr-bot

copy-pr-bot Bot commented Jul 3, 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.

@kashif

kashif commented Jul 3, 2026

Copy link
Copy Markdown
Contributor Author

One design question for reviewers on the attach_cp_kv_gather_hooks approach.

Right now the CP hook intercepts attention by swapping torch.nn.functional.scaled_dot_product_attention during each self_attn forward. That only works when the target lowers to torch SDPA, which is why cp_size>1 requires target_force_hf=true — a NeMo-custom target defaults to TE attention (and flex/tilelang/kernels-provided kernels would also bypass the swap), so the hook would silently no-op and each rank would fall back to local-only attention.

A more robust alternative: register a dedicated "cp_gather" entry in HF's ALL_ATTENTION_FUNCTIONS and select it via attn_implementation, instead of monkeypatching the torch symbol. That intercepts at the attention-interface level regardless of the underlying kernel, and would let us drop the force_hf requirement (and compose CP with faster attention backends). The trade-off is a bit more surface tied to the transformers attention-interface API across versions.

Worth doing in this PR, or fine as a follow-up? Leaning follow-up since the current gated version is correct where it's enabled, but happy to fold it in if preferred.

@kashif kashif marked this pull request as draft July 3, 2026 23:25
kashif added 7 commits July 3, 2026 23:40
…arallelism

Signed-off-by: Kashif Rasul <kashif.rasul@gmail.com>
Signed-off-by: Kashif Rasul <kashif.rasul@gmail.com>
Signed-off-by: Kashif Rasul <kashif.rasul@gmail.com>
…thout one

Signed-off-by: Kashif Rasul <kashif.rasul@gmail.com>
Signed-off-by: Kashif Rasul <kashif.rasul@gmail.com>
…fp32)

Signed-off-by: Kashif Rasul <kashif.rasul@gmail.com>
Signed-off-by: Kashif Rasul <kashif.rasul@gmail.com>
@khazic

khazic commented Jul 4, 2026

Copy link
Copy Markdown
Contributor

Taking a look at this now, will review shortly.

Comment thread nemo_automodel/recipes/llm/train_eagle3.py Outdated
Comment thread nemo_automodel/components/speculative/eagle/ring_attention.py
Comment thread nemo_automodel/recipes/llm/train_eagle3.py Outdated
Comment thread nemo_automodel/components/distributed/cp_utils.py Outdated
Comment thread nemo_automodel/components/speculative/dflash/target.py Outdated
@khazic

khazic commented Jul 4, 2026

Copy link
Copy Markdown
Contributor

One PR-level item: none of the 13 commits carries a Signed-off-by trailer, and CONTRIBUTING requires DCO sign-off on every commit. A rebase with --signoff will fix it.

kashif added 7 commits July 4, 2026 14:53
Signed-off-by: Kashif Rasul <kashif.rasul@gmail.com>
…rgets don't crash

Signed-off-by: Kashif Rasul <kashif.rasul@gmail.com>
Signed-off-by: Kashif Rasul <kashif.rasul@gmail.com>
…ngth breaks the cp preconditions

Signed-off-by: Kashif Rasul <kashif.rasul@gmail.com>
Signed-off-by: Kashif Rasul <kashif.rasul@gmail.com>
Signed-off-by: Kashif Rasul <kashif.rasul@gmail.com>
…raft trainers

Signed-off-by: Kashif Rasul <kashif.rasul@gmail.com>
@khazic

khazic commented Jul 4, 2026

Copy link
Copy Markdown
Contributor

There are a few more issues that I am still reviewing in depth.

Comment thread nemo_automodel/components/speculative/eagle/draft_llama.py Outdated
Comment thread nemo_automodel/recipes/llm/train_eagle3.py
Comment thread nemo_automodel/recipes/llm/train_eagle3.py
Comment thread nemo_automodel/recipes/llm/train_dspark.py
Comment thread nemo_automodel/components/distributed/cp_utils.py Outdated
Comment thread nemo_automodel/components/speculative/eagle/ring_attention.py
Comment thread nemo_automodel/components/speculative/eagle/ring_attention.py
Comment thread nemo_automodel/components/speculative/eagle/draft_llama.py Outdated
Comment thread nemo_automodel/components/speculative/eagle/ring_attention.py
Signed-off-by: Kashif Rasul <kashif.rasul@gmail.com>
@svcnvidia-nemo-ci svcnvidia-nemo-ci added the waiting-on-customer Waiting on the original author to respond label Jul 10, 2026
@HuiyingLi HuiyingLi merged commit 113d256 into NVIDIA-NeMo:main Jul 11, 2026
82 of 83 checks passed
@svcnvidia-nemo-ci svcnvidia-nemo-ci removed the waiting-on-customer Waiting on the original author to respond label Jul 11, 2026
@kashif kashif deleted the kashif/feat/draft-cp branch July 11, 2026 10:05
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.

4 participants