[Bug] CUDA graph backward replay illegal memory access after validation with mixed NVFP4/MXFP8 Mamba
Los mantenedores suelen responder en 2 días
Nadie ha tomado este issue todavía.
Evaluación
- Dificultad
- 5/5
- Tiempo estimado
- Más de una semana
- Aptitud para principiantes
- 35/100
- Tipo de issue
- Error
- Claridad
- Bastante claro
- Estado de actividad
- Tranquilo
- Área
- backend, machine-learning, performance
Línea de trabajo
Comienza con transformer_engine/pytorch/graph.py: backward y reproduce la transición de validación a entrenamiento de iteration-100 usando la configuración de GB200 indicada. Compara el Mamba CUDA graph que falla y la configuración mixta de NVFP4/MXFP8/BF16 con los controles que eliminan Mamba o usan E4M3 MXFP8; se considera terminado cuando el entrenamiento se reanuda después de la validación sin un illegal-memory-access ni un segmentation fault.
Escrito por el modelo de indexación a partir del texto del issue.
Descripción
Describe the bug
Nemotron 3 Super training on GB200 consistently hits an illegal memory access around the iteration-100 validation boundary when using Transformer Engine CUDA graphs with a mixed per-module NVFP4, MXFP8, and BF16 recipe.
- MXFP8 w/ cuda graphs for attn, mamba, moe_router, moe_preprocess: works
- NVFP4 w/ cuda graphs for attn, mamba, moe_router, moe_preprocess: hits illegal memory access between iterations 90 and 100
- NVFP4 w/ cuda graphs for attn, moe_router, moe_preprocess: works
The symptom initially appeared between the iteration 90 and 100 log points, but detailed timestamps show that iteration 100 and its 14 validation iterations both complete successfully. The failure occurs during the first training backward after validation. One rank reports a CUDA illegal-memory-access error from the Megatron Core backward call, while another rank can simultaneously segfault inside Transformer Engine CUDA graph backward replay. The failing ranks vary across executions.
Representative stacks:
train_step
-> forward_backward_no_pipelining
-> backward_step
-> custom_backward
-> Variable._execution_engine.run_backward
torch.AcceleratorError: CUDA error: an illegal memory access was encountered
transformer_engine/pytorch/graph.py: backward
-> torch/cuda/graphs.py: replay
-> at::cuda::CUDAGraph::replay
-> cudaGraphLaunch
-> cuGraphLaunch
Fatal Python error: Segmentation fault
Steps/Code to reproduce bug
Run Nemotron 3 Super training with this configuration:
Hardware: 16 GB200 nodes, 4 GPUs per node, 64 ranks
Parallelism: TP=2, PP=1, EP=64, ETP=1
Sequence length: 8192
Micro/global batch size: 1/512
Transformer implementation: transformer_engine
Attention backend: fused
CUDA graph implementation: transformer_engine
CUDA graph modules: [attn, mamba, moe_router, moe_preprocess]
CUDA graph warmup steps: 3
MoE dispatcher: flex with HybridEP, 32 SMs
Evaluation interval/iterations: 100/14
Manual GC interval: 100
The base quantization recipe is NVFP4 E2M1. The per-module precision configuration keeps attention QKV/projection, latent projections, and MTP in BF16, and uses MXFP8 for the Mamba output projection. The last 14 layers are also BF16.
The exact mixed-precision arguments used by the failing NVFP4 job were:
--bf16
--grad-reduce-in-bf16
--te-precision-config-file /mnt/artifacts/model/nemotron3_super_release_gb200/te_quant.cfg
--first-last-layers-bf16
--num-layers-at-start-in-bf16 0
--num-layers-at-end-in-bf16 14
--fp4-format e2m1
--fp4-recipe nvfp4
FP4 parameter gather was not enabled (fp4_param_gather=False).
The relevant runtime settings include:
CUDA_DEVICE_MAX_CONNECTIONS=32
NCCL_GRAPH_REGISTER=0
NCCL_NVLS_ENABLE=0
PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True
NUM_OF_HYBRID_EP_RANKS_PER_NVLINK_DOMAIN=64
NUM_OF_TOKENS_PER_CHUNK_COMBINE_API=128
NVTE_FWD_LAYERNORM_SM_MARGIN=20
NVTE_BWD_LAYERNORM_SM_MARGIN=20
Train through iteration 100, run validation, and resume training. In failing runs, iteration 100 completes, validation takes about 112 seconds and completes, then the first resumed training backward faults within several seconds.
What we tested and how we narrowed it down
- The failure reproduced across multiple independent executions and retries on different nodes and ranks.
- It reproduced with Transformer Engine 2.18 at two revisions, including
2.18.0+e0a871e6with the merged weight-preswizzling stale-pointer fix and the later single-grouped-weight fixes present. Those fixes did not resolve this end-to-end failure. - It reproduced with both NVSHMEM 3.6.5 and 3.4.5, so the NVSHMEM upgrade is not the cause.
- We aligned CUDA, NCCL, allocator, layernorm-margin, HybridEP-domain, and token-chunk settings with known working GB200 recipes. The failure remained.
- A Transformer Engine 2.17 plus cuDNN 9.21 diagnostic completed only five iterations. It did not exercise the iteration-100 boundary, so it is not evidence for or against a version regression.
- Keeping the updated configuration but removing only Mamba from the CUDA graph module list allowed validation at iteration 100 to complete and training to continue cleanly through iteration 120, with zero skipped or NaN iterations.
- Keeping Mamba in the full CUDA graph module list but replacing the mixed per-module NVFP4/MXFP8/BF16 recipe with an E4M3 MXFP8 recipe plus BF16 boundary layers allowed validation at iteration 100 to complete and training to continue cleanly through iteration 170, again with zero skipped or NaN iterations.
These controls narrow the problem to an interaction that requires Mamba CUDA graph capture and the mixed per-module NVFP4/MXFP8/BF16 recipe, or state introduced by that recipe across the train-to-validation-to-train transition. It does not look like a generic HybridEP, NVSHMEM, or Mamba training failure.
Expected behavior
Validation should not leave stale or incompatible quantization or CUDA graph state. Training should resume after validation without an illegal memory access during backward graph replay.
Environment overview
- Environment location: Docker on an internal bare-metal GB200 cluster
- Transformer Engine install: built from source from the Megatron Core dependency lock during the Docker image build
- Base image: NVIDIA PyTorch 26.06-derived image
Environment details
- Python: 3.12
- PyTorch:
2.13.0a0+8145d630e8.nv26.6 - Megatron Core:
0.19.0+af9e4408d - Transformer Engine:
2.18.0+e0a871e6 - CUDA: 13.3
- NCCL:
2.30.7+cuda13.3
Device details
- GPU model: NVIDIA GB200
- Topology: 16 nodes, 4 GPUs per node
Additional context
The CUDA error is asynchronous, so the Python error site alone does not identify the captured kernel that first corrupts memory. The simultaneous native stack on another rank localizes the visible failure to Transformer Engine backward graph replay, but we have not yet isolated the exact Mamba or quantization kernel. We can run a launch-blocking diagnostic or help reduce this to a smaller reproducer if there is a preferred instrumentation path.
- Lenguaje dominante
- Python
- Estrellas
- 3.6k
- Forks
- 851
- Merge medio
- 4 d 20 h
- PR fusionados (30 d)
- 58
Preparar el entorno
- Sin Dockerfile ni archivo de Docker Compose
- Tiene una plantilla de pull request
- Leer la guía de contribución
Primeros pasos
- Lee el issue completo y luego la guía de contribución del proyecto.
- Comenta en el issue que vas a ocuparte — evita que dos personas hagan lo mismo.
- Haz un fork del repositorio y trabaja en una rama.
- Abre un pull request que haga referencia al número del issue.
Más de NVIDIA/TransformerEngine
-
[Bug] group_quantize_fp8_blockwise: mbarrier invalidated before other threads finish waiting on itAbierto
Dificultad 2/5 1-3 horas Aptitud para principiantes 82/100
NVIDIA/TransformerEngine#3647 ·
Los mantenedores suelen responder en 2 días
-
[PyTorch] fp8_cs_quantize fake implementation returns a vector inverse scale instead of a scalarPosiblemente ocupada @sanjana658 la tomó hace 5 días. Abierto
Dificultad 2/5 1-3 horas Aptitud para principiantes 82/100
NVIDIA/TransformerEngine#3636 · 2 comentarios ·
Los mantenedores suelen responder en 2 días
-
[Bug] Backend selection picks FA3 for training with head_dim_qk=192 / v_head_dim=128, but FA3 backward cannot run itQuizá libre de nuevo @modelpath-dev la tomó hace 39 días y no hay ningún pull request abierto. Abiertoattention
Dificultad 2/5 1-3 horas Aptitud para principiantes 85/100
NVIDIA/TransformerEngine#3481 · 4 comentarios ·
Los mantenedores suelen responder en 2 días
-
Increase MAX_TENSOR_NUMAbiertobug
Dificultad 2/5 1-3 horas Aptitud para principiantes 68/100
NVIDIA/TransformerEngine#2189 · 7 comentarios · 5 reacciones ·
Los mantenedores suelen responder en 2 días
-
Dificultad 4/5 3-5 días Aptitud para principiantes 48/100
NVIDIA/TransformerEngine#3668 ·
Los mantenedores suelen responder en 2 días
Todos los issues de NVIDIA/TransformerEngine
Issues similares
-
Dificultad 2/5 1-3 horas Aptitud para principiantes 78/100
aws-samples/sample-per-model-bedrock#191 ·
Los mantenedores suelen responder en 1 día
-
Docs viewer: inverted grey colormap, Phantom CT window -20 to 35 HU, smoothing up to 40 mmPosiblemente ocupada Un pull request vinculado a esta issue está abierto o ya se fusionó. Abiertoarea: docs enhancement
Dificultad 2/5 1-3 horas Aptitud para principiantes 64/100
PyTomography/PyTomography#331 ·
Los mantenedores suelen responder en 1 día
-
effort: low prio: low ready
Dificultad 2/5 1-3 horas Aptitud para principiantes 82/100
digitalfabrik/integreat-cms#4616 ·
Los mantenedores suelen responder en 5 días
-
Dificultad 2/5 1-3 horas Aptitud para principiantes 66/100
FuRongJun-1999/dsh-memory#226 · 1 comentario ·
Los mantenedores suelen responder en 1 día
-
bug difficulty:2 good first issue help wanted
Dificultad 2/5 1-3 horas Aptitud para principiantes 88/100
MakazhanAlpamys/Soup#1865 ·
Los mantenedores suelen responder en 1 día