deepseek-ai/DeepEP
1
0
Fork 0
mirror of https://github.com/deepseek-ai/DeepEP.git synced 2026-11-05 00:47:32 +00:00
Code Issues 308 Projects Packages Wiki Activity

[GH-ISSUE #49] Assertion failed: internode_ll.cu:259, condition: num_recv_tokens != 0 #46

New issue
Open
opened 2026-10-26 09:53:12 +00:00 by bsakaguchi · 5 comments
bsakaguchi commented 2026-10-26 09:53:12 +00:00
Owner
Copy link

Originally created by @abcdabcd987 on GitHub (Mar 5, 2025).
Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/49

Hi, I ran into the following assertion when running the low latency kernel test over 4 nodes.

$ RANK=0 MASTER_ADDR=192.168.0.117 WORLD_SIZE=4 python tests/test_low_latency.py                                                                  
Allocating buffer size: 2116.290944 MB ...                                                                                                                                                                         
Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0                                                                                                                 
Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0                                                                                                                 
Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0                                                                                                                 
Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0                                                                                                                 
Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0                                                                                                                 
Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0                                                                                                                 
Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0                                                                                                                 
[rank5]:[E305 23:13:27.395915667 ProcessGroupNCCL.cpp:1895] [PG ID 1 PG GUID 1 Rank 5] Process group watchdog thread terminated with exception: CUDA error: unspecified launch failure                             
CUDA kernel errors might be asynchronously reported at some other API call, so the stacktrace below might be incorrect.                                                                                            
For debugging consider passing CUDA_LAUNCH_BLOCKING=1                                                                                                                                                              
Compile with `TORCH_USE_CUDA_DSA` to enable device-side assertions.                                                                                                                                                
                                                                                                                                                                                                                   
Exception raised from c10_cuda_check_implementation at /pytorch/c10/cuda/CUDAException.cpp:43 (most recent call first):                                                                                            
frame #0: c10::Error::Error(c10::SourceLocation, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >) + 0x98 (0x7f47afd7d788 in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libc10.so)                                                                                                                                                                                               
frame #1: c10::detail::torchCheckFail(char const*, char const*, unsigned int, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&) + 0xe0 (0x7f47afd26ea8 in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libc10.so)                                                                                                                                                           
frame #2: c10::cuda::c10_cuda_check_implementation(int, char const*, char const*, int, bool) + 0x3c2 (0x7f47b017e3d2 in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libc10_cuda.so)           
frame #3: c10d::ProcessGroupNCCL::WorkNCCL::finishedGPUExecutionInternal() const + 0x56 (0x7f475f5ce176 in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libtorch_cuda.so)                      
frame #4: c10d::ProcessGroupNCCL::WorkNCCL::isCompleted() + 0xa8 (0x7f475f5de658 in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libtorch_cuda.so)                                             
frame #5: c10d::ProcessGroupNCCL::watchdogHandler() + 0x63d (0x7f475f5dfbfd in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libtorch_cuda.so)                                                  
frame #6: c10d::ProcessGroupNCCL::ncclCommWatchdog() + 0x14d (0x7f475f5e0c4d in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libtorch_cuda.so)                                                 
frame #7: <unknown function> + 0xdc253 (0x7f474f8b0253 in /lib/x86_64-linux-gnu/libstdc++.so.6)                                                                                                                    
frame #8: <unknown function> + 0x94ac3 (0x7f47b0dc1ac3 in /lib/x86_64-linux-gnu/libc.so.6)                                                                                                                         
frame #9: <unknown function> + 0x126850 (0x7f47b0e53850 in /lib/x86_64-linux-gnu/libc.so.6)                                                                                                                        
                                                                                                                                                                                                                   
terminate called after throwing an instance of 'c10::DistBackendError'                                                                                                                                             
  what():  [PG ID 1 PG GUID 1 Rank 5] Process group watchdog thread terminated with exception: CUDA error: unspecified launch failure                                                                              
CUDA kernel errors might be asynchronously reported at some other API call, so the stacktrace below might be incorrect.                                                                                            
For debugging consider passing CUDA_LAUNCH_BLOCKING=1                                                                                                                                                              
Compile with `TORCH_USE_CUDA_DSA` to enable device-side assertions.

I tried to build the kernels with and without DISABLE_AGGRESSIVE_PTX_INSTRS=1. It did not help.

GPU is H100. NIC is ConnectX-7 (MT_0000000838). The link layer is Infiniband.

Error is the same even if it's on a single node (WORLD_SIZE=1).

Do you have any clue on where things might be wrong? Thanks!

Originally created by @abcdabcd987 on GitHub (Mar 5, 2025). Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/49 Hi, I ran into the following assertion when running the low latency kernel test over 4 nodes. ``` $ RANK=0 MASTER_ADDR=192.168.0.117 WORLD_SIZE=4 python tests/test_low_latency.py Allocating buffer size: 2116.290944 MB ... Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0 Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0 Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0 Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0 Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0 Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0 Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0 [rank5]:[E305 23:13:27.395915667 ProcessGroupNCCL.cpp:1895] [PG ID 1 PG GUID 1 Rank 5] Process group watchdog thread terminated with exception: CUDA error: unspecified launch failure CUDA kernel errors might be asynchronously reported at some other API call, so the stacktrace below might be incorrect. For debugging consider passing CUDA_LAUNCH_BLOCKING=1 Compile with `TORCH_USE_CUDA_DSA` to enable device-side assertions. Exception raised from c10_cuda_check_implementation at /pytorch/c10/cuda/CUDAException.cpp:43 (most recent call first): frame #0: c10::Error::Error(c10::SourceLocation, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >) + 0x98 (0x7f47afd7d788 in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libc10.so) frame #1: c10::detail::torchCheckFail(char const*, char const*, unsigned int, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&) + 0xe0 (0x7f47afd26ea8 in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libc10.so) frame #2: c10::cuda::c10_cuda_check_implementation(int, char const*, char const*, int, bool) + 0x3c2 (0x7f47b017e3d2 in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libc10_cuda.so) frame #3: c10d::ProcessGroupNCCL::WorkNCCL::finishedGPUExecutionInternal() const + 0x56 (0x7f475f5ce176 in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libtorch_cuda.so) frame #4: c10d::ProcessGroupNCCL::WorkNCCL::isCompleted() + 0xa8 (0x7f475f5de658 in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libtorch_cuda.so) frame #5: c10d::ProcessGroupNCCL::watchdogHandler() + 0x63d (0x7f475f5dfbfd in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libtorch_cuda.so) frame #6: c10d::ProcessGroupNCCL::ncclCommWatchdog() + 0x14d (0x7f475f5e0c4d in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libtorch_cuda.so) frame #7: <unknown function> + 0xdc253 (0x7f474f8b0253 in /lib/x86_64-linux-gnu/libstdc++.so.6) frame #8: <unknown function> + 0x94ac3 (0x7f47b0dc1ac3 in /lib/x86_64-linux-gnu/libc.so.6) frame #9: <unknown function> + 0x126850 (0x7f47b0e53850 in /lib/x86_64-linux-gnu/libc.so.6) terminate called after throwing an instance of 'c10::DistBackendError' what(): [PG ID 1 PG GUID 1 Rank 5] Process group watchdog thread terminated with exception: CUDA error: unspecified launch failure CUDA kernel errors might be asynchronously reported at some other API call, so the stacktrace below might be incorrect. For debugging consider passing CUDA_LAUNCH_BLOCKING=1 Compile with `TORCH_USE_CUDA_DSA` to enable device-side assertions. ``` I tried to build the kernels with and without `DISABLE_AGGRESSIVE_PTX_INSTRS=1`. It did not help. GPU is H100. NIC is ConnectX-7 (MT_0000000838). The link layer is Infiniband. Error is the same even if it's on a single node (`WORLD_SIZE=1`). Do you have any clue on where things might be wrong? Thanks!
bsakaguchi commented 2026-10-26 09:53:13 +00:00
Author
Owner
Copy link

@abcdabcd987 commented on GitHub (Mar 5, 2025):

test_low_latency.txt

Here's the full log with NVSHMEM_DEBUG=TRACE if it helps.

<!-- gh-comment-id:2702317674 --> @abcdabcd987 commented on GitHub (Mar 5, 2025): [test_low_latency.txt](https://github.com/user-attachments/files/19098395/test_low_latency.txt) Here's the full log with `NVSHMEM_DEBUG=TRACE` if it helps.
bsakaguchi commented 2026-10-26 09:53:13 +00:00
Author
Owner
Copy link

@LyricZhao commented on GitHub (Mar 6, 2025):

Can you confirm the NVSHMEM patch is correctly installed? This seems something wrong with network adaptive routing (AR). Or you can try switch on/off the AR setting for your network.

cc @sphish

<!-- gh-comment-id:2702504824 --> @LyricZhao commented on GitHub (Mar 6, 2025): Can you confirm the NVSHMEM patch is correctly installed? This seems something wrong with network adaptive routing (AR). Or you can try switch on/off the AR setting for your network. cc @sphish
bsakaguchi commented 2026-10-26 09:53:13 +00:00
Author
Owner
Copy link

@LyricZhao commented on GitHub (Mar 6, 2025):

Or could you please try modifying num_recv_tokens = ld_acquire_global(rdma_recv_count + local_expert_idx * num_ranks + src_rank); into ld_acquire_sys_global?

<!-- gh-comment-id:2702532731 --> @LyricZhao commented on GitHub (Mar 6, 2025): Or could you please try modifying `num_recv_tokens = ld_acquire_global(rdma_recv_count + local_expert_idx * num_ranks + src_rank);` into `ld_acquire_sys_global`?
bsakaguchi commented 2026-10-26 09:53:14 +00:00
Author
Owner
Copy link

@GangLuICT commented on GitHub (Mar 13, 2025):

Hi, I ran into the following assertion when running the low latency kernel test over 4 nodes.

$ RANK=0 MASTER_ADDR=192.168.0.117 WORLD_SIZE=4 python tests/test_low_latency.py                                                                  
Allocating buffer size: 2116.290944 MB ...                                                                                                                                                                         
Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0                                                                                                                 
Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0                                                                                                                 
Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0                                                                                                                 
Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0                                                                                                                 
Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0                                                                                                                 
Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0                                                                                                                 
Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0                                                                                                                 
[rank5]:[E305 23:13:27.395915667 ProcessGroupNCCL.cpp:1895] [PG ID 1 PG GUID 1 Rank 5] Process group watchdog thread terminated with exception: CUDA error: unspecified launch failure                             
CUDA kernel errors might be asynchronously reported at some other API call, so the stacktrace below might be incorrect.                                                                                            
For debugging consider passing CUDA_LAUNCH_BLOCKING=1                                                                                                                                                              
Compile with `TORCH_USE_CUDA_DSA` to enable device-side assertions.                                                                                                                                                
                                                                                                                                                                                                                   
Exception raised from c10_cuda_check_implementation at /pytorch/c10/cuda/CUDAException.cpp:43 (most recent call first):                                                                                            
frame #0: c10::Error::Error(c10::SourceLocation, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >) + 0x98 (0x7f47afd7d788 in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libc10.so)                                                                                                                                                                                               
frame #1: c10::detail::torchCheckFail(char const*, char const*, unsigned int, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&) + 0xe0 (0x7f47afd26ea8 in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libc10.so)                                                                                                                                                           
frame #2: c10::cuda::c10_cuda_check_implementation(int, char const*, char const*, int, bool) + 0x3c2 (0x7f47b017e3d2 in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libc10_cuda.so)           
frame #3: c10d::ProcessGroupNCCL::WorkNCCL::finishedGPUExecutionInternal() const + 0x56 (0x7f475f5ce176 in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libtorch_cuda.so)                      
frame #4: c10d::ProcessGroupNCCL::WorkNCCL::isCompleted() + 0xa8 (0x7f475f5de658 in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libtorch_cuda.so)                                             
frame #5: c10d::ProcessGroupNCCL::watchdogHandler() + 0x63d (0x7f475f5dfbfd in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libtorch_cuda.so)                                                  
frame #6: c10d::ProcessGroupNCCL::ncclCommWatchdog() + 0x14d (0x7f475f5e0c4d in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libtorch_cuda.so)                                                 
frame #7: <unknown function> + 0xdc253 (0x7f474f8b0253 in /lib/x86_64-linux-gnu/libstdc++.so.6)                                                                                                                    
frame #8: <unknown function> + 0x94ac3 (0x7f47b0dc1ac3 in /lib/x86_64-linux-gnu/libc.so.6)                                                                                                                         
frame #9: <unknown function> + 0x126850 (0x7f47b0e53850 in /lib/x86_64-linux-gnu/libc.so.6)                                                                                                                        
                                                                                                                                                                                                                   
terminate called after throwing an instance of 'c10::DistBackendError'                                                                                                                                             
  what():  [PG ID 1 PG GUID 1 Rank 5] Process group watchdog thread terminated with exception: CUDA error: unspecified launch failure                                                                              
CUDA kernel errors might be asynchronously reported at some other API call, so the stacktrace below might be incorrect.                                                                                            
For debugging consider passing CUDA_LAUNCH_BLOCKING=1                                                                                                                                                              
Compile with `TORCH_USE_CUDA_DSA` to enable device-side assertions.

I tried to build the kernels with and without DISABLE_AGGRESSIVE_PTX_INSTRS=1. It did not help.

GPU is H100. NIC is ConnectX-7 (MT_0000000838). The link layer is Infiniband.

Error is the same even if it's on a single node (WORLD_SIZE=1).

Do you have any clue on where things might be wrong? Thanks!

Have you solved this issue? We met the same issue on a SRIOV/VF & iommu=on environment.
But it does not occur when SRIOV is disabled.

<!-- gh-comment-id:2720712225 --> @GangLuICT commented on GitHub (Mar 13, 2025): > Hi, I ran into the following assertion when running the low latency kernel test over 4 nodes. > > ``` > $ RANK=0 MASTER_ADDR=192.168.0.117 WORLD_SIZE=4 python tests/test_low_latency.py > Allocating buffer size: 2116.290944 MB ... > Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0 > Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0 > Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0 > Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0 > Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0 > Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0 > Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0 > [rank5]:[E305 23:13:27.395915667 ProcessGroupNCCL.cpp:1895] [PG ID 1 PG GUID 1 Rank 5] Process group watchdog thread terminated with exception: CUDA error: unspecified launch failure > CUDA kernel errors might be asynchronously reported at some other API call, so the stacktrace below might be incorrect. > For debugging consider passing CUDA_LAUNCH_BLOCKING=1 > Compile with `TORCH_USE_CUDA_DSA` to enable device-side assertions. > > Exception raised from c10_cuda_check_implementation at /pytorch/c10/cuda/CUDAException.cpp:43 (most recent call first): > frame #0: c10::Error::Error(c10::SourceLocation, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >) + 0x98 (0x7f47afd7d788 in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libc10.so) > frame #1: c10::detail::torchCheckFail(char const*, char const*, unsigned int, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&) + 0xe0 (0x7f47afd26ea8 in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libc10.so) > frame #2: c10::cuda::c10_cuda_check_implementation(int, char const*, char const*, int, bool) + 0x3c2 (0x7f47b017e3d2 in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libc10_cuda.so) > frame #3: c10d::ProcessGroupNCCL::WorkNCCL::finishedGPUExecutionInternal() const + 0x56 (0x7f475f5ce176 in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libtorch_cuda.so) > frame #4: c10d::ProcessGroupNCCL::WorkNCCL::isCompleted() + 0xa8 (0x7f475f5de658 in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libtorch_cuda.so) > frame #5: c10d::ProcessGroupNCCL::watchdogHandler() + 0x63d (0x7f475f5dfbfd in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libtorch_cuda.so) > frame #6: c10d::ProcessGroupNCCL::ncclCommWatchdog() + 0x14d (0x7f475f5e0c4d in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libtorch_cuda.so) > frame #7: <unknown function> + 0xdc253 (0x7f474f8b0253 in /lib/x86_64-linux-gnu/libstdc++.so.6) > frame #8: <unknown function> + 0x94ac3 (0x7f47b0dc1ac3 in /lib/x86_64-linux-gnu/libc.so.6) > frame #9: <unknown function> + 0x126850 (0x7f47b0e53850 in /lib/x86_64-linux-gnu/libc.so.6) > > terminate called after throwing an instance of 'c10::DistBackendError' > what(): [PG ID 1 PG GUID 1 Rank 5] Process group watchdog thread terminated with exception: CUDA error: unspecified launch failure > CUDA kernel errors might be asynchronously reported at some other API call, so the stacktrace below might be incorrect. > For debugging consider passing CUDA_LAUNCH_BLOCKING=1 > Compile with `TORCH_USE_CUDA_DSA` to enable device-side assertions. > ``` > > I tried to build the kernels with and without `DISABLE_AGGRESSIVE_PTX_INSTRS=1`. It did not help. > > GPU is H100. NIC is ConnectX-7 (MT_0000000838). The link layer is Infiniband. > > Error is the same even if it's on a single node (`WORLD_SIZE=1`). > > Do you have any clue on where things might be wrong? Thanks! Have you solved this issue? We met the same issue on a SRIOV/VF & iommu=on environment. But it does not occur when SRIOV is disabled.
bsakaguchi commented 2026-10-26 09:53:14 +00:00
Author
Owner
Copy link

@sphish commented on GitHub (Mar 18, 2025):

Hi, I ran into the following assertion when running the low latency kernel test over 4 nodes.

$ RANK=0 MASTER_ADDR=192.168.0.117 WORLD_SIZE=4 python tests/test_low_latency.py                                                                  
Allocating buffer size: 2116.290944 MB ...                                                                                                                                                                         
Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0                                                                                                                 
Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0                                                                                                                 
Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0                                                                                                                 
Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0                                                                                                                 
Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0                                                                                                                 
Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0                                                                                                                 
Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0                                                                                                                 
[rank5]:[E305 23:13:27.395915667 ProcessGroupNCCL.cpp:1895] [PG ID 1 PG GUID 1 Rank 5] Process group watchdog thread terminated with exception: CUDA error: unspecified launch failure                             
CUDA kernel errors might be asynchronously reported at some other API call, so the stacktrace below might be incorrect.                                                                                            
For debugging consider passing CUDA_LAUNCH_BLOCKING=1                                                                                                                                                              
Compile with `TORCH_USE_CUDA_DSA` to enable device-side assertions.                                                                                                                                                
                                                                                                                                                                                                                   
Exception raised from c10_cuda_check_implementation at /pytorch/c10/cuda/CUDAException.cpp:43 (most recent call first):                                                                                            
frame #0: c10::Error::Error(c10::SourceLocation, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >) + 0x98 (0x7f47afd7d788 in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libc10.so)                                                                                                                                                                                               
frame #1: c10::detail::torchCheckFail(char const*, char const*, unsigned int, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&) + 0xe0 (0x7f47afd26ea8 in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libc10.so)                                                                                                                                                           
frame #2: c10::cuda::c10_cuda_check_implementation(int, char const*, char const*, int, bool) + 0x3c2 (0x7f47b017e3d2 in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libc10_cuda.so)           
frame #3: c10d::ProcessGroupNCCL::WorkNCCL::finishedGPUExecutionInternal() const + 0x56 (0x7f475f5ce176 in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libtorch_cuda.so)                      
frame #4: c10d::ProcessGroupNCCL::WorkNCCL::isCompleted() + 0xa8 (0x7f475f5de658 in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libtorch_cuda.so)                                             
frame #5: c10d::ProcessGroupNCCL::watchdogHandler() + 0x63d (0x7f475f5dfbfd in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libtorch_cuda.so)                                                  
frame #6: c10d::ProcessGroupNCCL::ncclCommWatchdog() + 0x14d (0x7f475f5e0c4d in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libtorch_cuda.so)                                                 
frame #7: <unknown function> + 0xdc253 (0x7f474f8b0253 in /lib/x86_64-linux-gnu/libstdc++.so.6)                                                                                                                    
frame #8: <unknown function> + 0x94ac3 (0x7f47b0dc1ac3 in /lib/x86_64-linux-gnu/libc.so.6)                                                                                                                         
frame #9: <unknown function> + 0x126850 (0x7f47b0e53850 in /lib/x86_64-linux-gnu/libc.so.6)                                                                                                                        
                                                                                                                                                                                                                   
terminate called after throwing an instance of 'c10::DistBackendError'                                                                                                                                             
  what():  [PG ID 1 PG GUID 1 Rank 5] Process group watchdog thread terminated with exception: CUDA error: unspecified launch failure                                                                              
CUDA kernel errors might be asynchronously reported at some other API call, so the stacktrace below might be incorrect.                                                                                            
For debugging consider passing CUDA_LAUNCH_BLOCKING=1                                                                                                                                                              
Compile with `TORCH_USE_CUDA_DSA` to enable device-side assertions.

I tried to build the kernels with and without DISABLE_AGGRESSIVE_PTX_INSTRS=1. It did not help.
GPU is H100. NIC is ConnectX-7 (MT_0000000838). The link layer is Infiniband.
Error is the same even if it's on a single node (WORLD_SIZE=1).
Do you have any clue on where things might be wrong? Thanks!

Have you solved this issue? We met the same issue on a SRIOV/VF & iommu=on environment. But it does not occur when SRIOV is disabled.

It seems that one of the ranks might have encountered an exception and didn't reach the GPU kernel execution. We have also disabled SRIOV and haven't experienced similar issues before.

<!-- gh-comment-id:2732245208 --> @sphish commented on GitHub (Mar 18, 2025): > > Hi, I ran into the following assertion when running the low latency kernel test over 4 nodes. > > ``` > > $ RANK=0 MASTER_ADDR=192.168.0.117 WORLD_SIZE=4 python tests/test_low_latency.py > > Allocating buffer size: 2116.290944 MB ... > > Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0 > > Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0 > > Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0 > > Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0 > > Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0 > > Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0 > > Assertion failed: /mnt/fs/DeepEP/csrc/kernels/internode_ll.cu:259, condition: num_recv_tokens != 0 > > [rank5]:[E305 23:13:27.395915667 ProcessGroupNCCL.cpp:1895] [PG ID 1 PG GUID 1 Rank 5] Process group watchdog thread terminated with exception: CUDA error: unspecified launch failure > > CUDA kernel errors might be asynchronously reported at some other API call, so the stacktrace below might be incorrect. > > For debugging consider passing CUDA_LAUNCH_BLOCKING=1 > > Compile with `TORCH_USE_CUDA_DSA` to enable device-side assertions. > > > > Exception raised from c10_cuda_check_implementation at /pytorch/c10/cuda/CUDAException.cpp:43 (most recent call first): > > frame #0: c10::Error::Error(c10::SourceLocation, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >) + 0x98 (0x7f47afd7d788 in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libc10.so) > > frame #1: c10::detail::torchCheckFail(char const*, char const*, unsigned int, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&) + 0xe0 (0x7f47afd26ea8 in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libc10.so) > > frame #2: c10::cuda::c10_cuda_check_implementation(int, char const*, char const*, int, bool) + 0x3c2 (0x7f47b017e3d2 in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libc10_cuda.so) > > frame #3: c10d::ProcessGroupNCCL::WorkNCCL::finishedGPUExecutionInternal() const + 0x56 (0x7f475f5ce176 in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libtorch_cuda.so) > > frame #4: c10d::ProcessGroupNCCL::WorkNCCL::isCompleted() + 0xa8 (0x7f475f5de658 in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libtorch_cuda.so) > > frame #5: c10d::ProcessGroupNCCL::watchdogHandler() + 0x63d (0x7f475f5dfbfd in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libtorch_cuda.so) > > frame #6: c10d::ProcessGroupNCCL::ncclCommWatchdog() + 0x14d (0x7f475f5e0c4d in /mnt/fs/conda/envs/DeepEP/lib/python3.10/site-packages/torch/lib/libtorch_cuda.so) > > frame #7: <unknown function> + 0xdc253 (0x7f474f8b0253 in /lib/x86_64-linux-gnu/libstdc++.so.6) > > frame #8: <unknown function> + 0x94ac3 (0x7f47b0dc1ac3 in /lib/x86_64-linux-gnu/libc.so.6) > > frame #9: <unknown function> + 0x126850 (0x7f47b0e53850 in /lib/x86_64-linux-gnu/libc.so.6) > > > > terminate called after throwing an instance of 'c10::DistBackendError' > > what(): [PG ID 1 PG GUID 1 Rank 5] Process group watchdog thread terminated with exception: CUDA error: unspecified launch failure > > CUDA kernel errors might be asynchronously reported at some other API call, so the stacktrace below might be incorrect. > > For debugging consider passing CUDA_LAUNCH_BLOCKING=1 > > Compile with `TORCH_USE_CUDA_DSA` to enable device-side assertions. > > ``` > > > > > > > > > > I tried to build the kernels with and without `DISABLE_AGGRESSIVE_PTX_INSTRS=1`. It did not help. > > GPU is H100. NIC is ConnectX-7 (MT_0000000838). The link layer is Infiniband. > > Error is the same even if it's on a single node (`WORLD_SIZE=1`). > > Do you have any clue on where things might be wrong? Thanks! > > Have you solved this issue? We met the same issue on a SRIOV/VF & iommu=on environment. But it does not occur when SRIOV is disabled. It seems that one of the ranks might have encountered an exception and didn't reach the GPU kernel execution. We have also disabled SRIOV and haven't experienced similar issues before.
Sign in to join this conversation.
No Branch/Tag specified
Branches Tags
main
hybrid-ep
chore/remove-clang-format-yapf
v2.5
nvDev
try-fix-cu128
mori-ep
antgroup-opt
compat-old-torch
add-linter
eager-rdma
v1.2.1
Labels
Clear labels   
pull-request
Mirrored from GitHub Pull Request

No labels
pull-request
Milestone
Clear milestone
No items
No milestone
Projects
Clear projects
No items
No project
Assignees
Clear assignees
No assignees
1 participant
Notifications
Due date
The due date is invalid or out of range. Please use the format "yyyy-mm-dd".

No due date set.

Dependencies

No dependencies set.

Reference
deepseek-ai/DeepEP#46
No description provided.