[Bug]: Find the root cause of SHARED_EXPERTS_STREAM fail

[vadiklyutiy]

In #28180 there was implemented workaround for apply_repetition_penalties fail.

Need to investigate the root cause.

All preliminary details in #28180

[vadiklyutiy]

I found if set cuda.torch.compile() after self.model_executor.execute_model we got fail on this sync.

�[1;36m(EngineCore_DP0 pid=673114)�[0;0m ERROR 11-07 05:05:01 [core.py:846] EngineCore encountered a fatal error.
�[1;36m(EngineCore_DP0 pid=673114)�[0;0m ERROR 11-07 05:05:01 [core.py:846] Traceback (most recent call last):
�[1;36m(EngineCore_DP0 pid=673114)�[0;0m ERROR 11-07 05:05:01 [core.py:846]   File "/home/scratch.vgimpelson_ent/vllm_qwen2/vllm/v1/engine/core.py", line 837, in run_engine_core
�[1;36m(EngineCore_DP0 pid=673114)�[0;0m ERROR 11-07 05:05:01 [core.py:846]     engine_core.run_busy_loop()
�[1;36m(EngineCore_DP0 pid=673114)�[0;0m ERROR 11-07 05:05:01 [core.py:846]   File "/home/scratch.vgimpelson_ent/vllm_qwen2/vllm/v1/engine/core.py", line 864, in run_busy_loop
�[1;36m(EngineCore_DP0 pid=673114)�[0;0m ERROR 11-07 05:05:01 [core.py:846]     self._process_engine_step()
�[1;36m(EngineCore_DP0 pid=673114)�[0;0m ERROR 11-07 05:05:01 [core.py:846]   File "/home/scratch.vgimpelson_ent/vllm_qwen2/vllm/v1/engine/core.py", line 893, in _process_engine_step
�[1;36m(EngineCore_DP0 pid=673114)�[0;0m ERROR 11-07 05:05:01 [core.py:846]     outputs, model_executed = self.step_fn()
�[1;36m(EngineCore_DP0 pid=673114)�[0;0m ERROR 11-07 05:05:01 [core.py:846]                               ^^^^^^^^^^^^^^
�[1;36m(EngineCore_DP0 pid=673114)�[0;0m ERROR 11-07 05:05:01 [core.py:846]   File "/home/scratch.vgimpelson_ent/vllm_qwen2/vllm/v1/engine/core.py", line 338, in step
�[1;36m(EngineCore_DP0 pid=673114)�[0;0m ERROR 11-07 05:05:01 [core.py:846]     torch.cuda.synchronize()
�[1;36m(EngineCore_DP0 pid=673114)�[0;0m ERROR 11-07 05:05:01 [core.py:846]   File "/home/scratch.vgimpelson_ent/venv_l4/lib/python3.12/site-packages/torch/cuda/__init__.py", line 1083, in synchronize
�[1;36m(EngineCore_DP0 pid=673114)�[0;0m ERROR 11-07 05:05:01 [core.py:846]     return torch._C._cuda_synchronize()
�[1;36m(EngineCore_DP0 pid=673114)�[0;0m ERROR 11-07 05:05:01 [core.py:846]            ^^^^^^^^^^^^^^^^^^^^^^^^^^^^
�[1;36m(EngineCore_DP0 pid=673114)�[0;0m ERROR 11-07 05:05:01 [core.py:846] torch.AcceleratorError: CUDA error: an illegal memory access was encountered
�[1;36m(EngineCore_DP0 pid=673114)�[0;0m ERROR 11-07 05:05:01 [core.py:846] Search for `cudaErrorIllegalAddress' in https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__TYPES.html for more information.
�[1;36m(EngineCore_DP0 pid=673114)�[0;0m ERROR 11-07 05:05:01 [core.py:846] CUDA kernel errors might be asynchronously reported at some other API call, so the stacktrace below might be incorrect.
�[1;36m(EngineCore_DP0 pid=673114)�[0;0m ERROR 11-07 05:05:01 [core.py:846] For debugging consider passing CUDA_LAUNCH_BLOCKING=1
�[1;36m(EngineCore_DP0 pid=673114)�[0;0m ERROR 11-07 05:05:01 [core.py:846] Compile with `TORCH_USE_CUDA_DSA` to enable device-side assertions.

Further I was following this guide https://blog.vllm.ai/2025/08/11/cuda-debugging.html (thanks @youkaichao for pointing me there) and collected coredump

export CUDA_ENABLE_COREDUMP_ON_EXCEPTION=1 CUDA_COREDUMP_SHOW_PROGRESS=1 CUDA_COREDUMP_GENERATION_FLAGS='skip_nonrelocated_elf_images,skip_global_memory,skip_shared_memory,skip_local_memory,skip_constbank_memory' CUDA_COREDUMP_FILE="/tmp/cuda_coredump_%h.%p.%t"

for i in {1..10}; do pytest -sv tests/evals/gsm8k/test_gsm8k_correctness.py::test_gsm8k_correctness_param[Qwen1.5-MoE-W4A16-CT-tp1]; done

(need loop because bug is not regular)

cuda-gdb show the following

(cuda-gdb) target cudacore /tmp/cuda_coredump_ipp2-0566.708702.1762510498
Opening GPU coredump: /tmp/cuda_coredump_ipp2-0566.708702.1762510498
[Current focus set to CUDA kernel 0, grid 19500, block (52,0,0), thread (32,0,0), device 0, sm 33, warp 0, lane 0]

CUDA Exception: Warp Illegal Address
The exception was triggered at PC 0x7cbdc74cd9b0  void marlin_moe_wna16::Marlin<__nv_bfloat16, 1125899907892224l, 1125899906909960l, 256, 4, 16, 4, false, 4, 8, false>(int4 const*, int4 const*, int4*, int4*, int4 const*, int4 const*, unsigned short const*, int4 const*, int const*, int const*, int const*, int const*, float const*, int, bool, bool, int, int, int, int, int*, bool, bool, bool, int)
#0  0x00007cbdc74cd9f0 in void marlin_moe_wna16::Marlin<__nv_bfloat16, 1125899907892224l, 1125899906909960l, 256, 4, 16, 4, false, 4, 8, false>(int4 const*, int4 const*, int4*, int4*, int4 const*, int4 const*, unsigned short const*, int4 const*, int const*, int const*, int const*, int const*, float const*, int, bool, bool, int, int, int, int, int*, bool, bool, bool, int)<<<(58,1,1),(256,1,1)>>> ()

cc @simon-mo

[zhewenl]

Temp fix to CI: #28324
We need to investigate the root cause for a forward fix to Qwen1.5-MoE-W4A16-CT-tp1 and other MoE models affected by this

[jinzhen-lin]

@vadiklyutiy Which gpu are you used ? I found that the launch setting of marlin kernel is (58,1,1), which means that 58 SMs are used (in marlin kernel, num of threadblocks = num os SMs). I don't know which gpu has exactly 58 SMs.

[vadiklyutiy]

@vadiklyutiy Which gpu are you used ? I found that the launch setting of marlin kernel is (58,1,1), which means that 58 SMs are used (in marlin kernel, num of threadblocks = num os SMs). I don't know which gpu has exactly 58 SMs.

I use L4 that has 58 SMs

$ python -c "import torch; print(torch.cuda.get_device_properties(0).multi_processor_count)"
58