feat(metax): FSDP2 full feature parity + Qwen3 training match - #62
Merged
Conversation
Enable FSDP2 (fully_shard + DTensor) on MetaX boxing with complete feature coverage: per-layer wrapping, mixed-precision, optimizers, gradient clipping, state_dict, no-sync accumulation, CPU offload, and 2D mesh (DP×TP). Verified against Qwen3-0.6B training: 20-step trajectory matches native MetaX torch to 7.6e-08 relative diff (3.17→0.034 loss on both), throughput 0.98×. Root cause: blocking cudaMemcpy only orders against the legacy default stream, so FSDP2's reduce-scatter-stream D2H copied stale data (zeros). Fixed by draining the current stream before every blocking Memcpy in copy_ops.cc. Changes: - copy_ops.cc: SyncCurrentStreamBeforeBlockingCopy() at 7 call sites (after each contiguous() that may enqueue work). Guards _to_copy CPU-offload, _copy_from cross-device, and _local_scalar_dense. - codegen_ops.py: box plain const Tensor& args in foreach ops (37 kernels) to prevent self-recursion SIGSEGV when unboxed inputs re-dispatch. - process_group.py: _register_inner_backend() so DeviceMesh can read pg.group_name (ProcessGroup::setGroupName forwards only to registered). - test_fsdp2_features_metax.py: 13 checks, ws4, all pass. - test_qwen3_fsdp2_metax.py: honest flagos-vs-native comparison (no torch_fl import in native mode, identical checkpoint+batches, pad-masked labels). - test_fsdp_live_metax.py: FSDP1+FSDP2 smoke tests (9 checks). Known upstream regression (NOT caused by this work, isolated by bisection): FSDP1 at world_size≥4 is flaky on 11457c9 (profiler commit that enabled real CUDA streams in guard.h). ws4 ~50% diverge at step 1, ws8 produces garbage. My work is 5/5 stable on both 8695f5a and d530681; pure upstream SIGSEGVs without the codegen fix. FSDP2 is unaffected. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
zhaoyinglia
approved these changes
Aug 6, 2026
zhaoyinglia
approved these changes
Aug 6, 2026
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.
Enable FSDP2 (
fully_shard+ DTensor) on MetaX boxing with complete feature coverage: per-layer wrapping, mixed-precision, optimizers, gradient clipping, state_dict, no-sync accumulation, CPU offload, and 2D mesh (DP×TP).Verification
FSDP2 features (13/13 pass, ws4):
clip_grad_norm_→ matches single-GPUQwen3-0.6B training (20 steps, ws4, batch 2, seq 128):
Ops integration: 589 passed, 259 skipped, 1 xfailed, 3 xpassed (9m4s)
Root Cause
Blocking
cudaMemcpy(used by FSDP2 CPU offload) only orders against the legacy default stream. FSDP2's reduce-scatter runs in a side stream, so the D2H copy saw zeros. PyTorch creates side streams withcudaStreamNonBlocking, making this deterministic wrong-data (not a race).Changes
csrc/aten/copy_ops.cc(+52 lines):SyncCurrentStreamBeforeBlockingCopy()at 7 call sites — after eachcontiguous()that may enqueue work, before the blockingMemcpy. Guards_to_copyCPU-offload,_copy_fromcross-device, and_local_scalar_dense.scripts/codegen_ops.py(+37 lines): box plainconst Tensor&args in foreach ops (37 kernels) to prevent self-recursion SIGSEGV when unboxed inputs re-dispatch on PrivateUse1.torch_fl/comm/process_group.py(+30 lines):_register_inner_backend()soDeviceMeshcan readpg.group_name(ProcessGroup::setGroupNameforwards only to registered backends).Test files (+1078 lines):
test_fsdp2_features_metax.py: 13 checks covering all FSDP2 featurestest_qwen3_fsdp2_metax.py: honest flagos-vs-native comparison (notorch_flimport in native mode, identical checkpoint+batches, pad-masked labels)test_fsdp_live_metax.py: FSDP1+FSDP2 smoke tests (9 checks)Known Upstream Regression (not caused by this work)
FSDP1 at world_size≥4 is flaky on upstream
11457c9(profiler commit that enabled real CUDA streams inguard.h):Bisection confirms:
8695f5a→ 5/5 stabled530681→ 5/5 stable11457c9→ flakycopy_ops.cc→ still flaky (not mine)11457c9→ SIGSEGV (needs the codegen fix)FSDP2 is unaffected (13/13 features, Qwen3 match).
🤖 Generated with Claude Code