[GH-ISSUE #480] [bug]Missing some critical error log when timeout

bsakaguchi commented

2026-10-26 09:58:21 +00:00

Owner

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

1. Description

I encountered some errors just like this. Rdma_rank 1 have this error, the sender coordinator and receiver both have timeout error.

For NVL receicer, head == tail, it seem to be related to Forwarder.
For Sender coordinator, It depends on Sender to update rdma_send_channel_tail , so it seem to be related to Sender.

but both the forwarder and Sender does not, we can't know the forwarder or sender has error or not ?
I think it's related to trap(). After trap being called by on warp, The execution of the kernel is aborted, if other warp's print buffer can't be flushed to output, it may be lost.

DeepEP dispatch NVL receiver timeout, channel: 1, RDMA: 1, nvl: 1, src NVL: 2, head: 75, tail: 75
DeepEP RDMA sender coordinator timeout, channel: 1, IB: 1, nvl 2, dst IB: 0, tail: 128, remaining: 21
DeepEP RDMA sender coordinator timeout, channel: 1, IB: 1, nvl 2, dst IB: 1, tail: 128, remaining: 27
DeepEP RDMA sender coordinator timeout, channel: 1, IB: 1, nvl 2, dst IB: 2, tail: 128, remaining: 30
DeepEP RDMA sender coordinator timeout, channel: 1, IB: 1, nvl 2, dst IB: 3, tail: 112, remaining: 25
DeepEP dispatch NVL receiver timeout, channel: 1, RDMA: 1, nvl: 7, src NVL: 2, head: 85, tail: 85
DeepEP dispatch NVL receiver timeout, channel: 1, RDMA: 1, nvl: 6, src NVL: 2, head: 83, tail: 83
DeepEP dispatch NVL receiver timeout, channel: 1, RDMA: 1, nvl: 5, src NVL: 2, head: 88, tail: 88
DeepEP dispatch NVL receiver timeout, channel: 1, RDMA: 1, nvl: 0, src NVL: 2, head: 76, tail: 76
DeepEP dispatch NVL receiver timeout, channel: 1, RDMA: 1, nvl: 4, src NVL: 2, head: 86, tail: 86
DeepEP dispatch NVL receiver timeout, channel: 1, RDMA: 1, nvl: 3, src NVL: 2, head: 84, tail: 84

2. Verification

Just to verify the phenomenon exist, the program can't output the all warp role, we can only see few warp output in same SM. I add some sleep before trap, it seem all warp have output, But it may not be guaranteed.
So is there a way to ensure that all warp outputs are guaranteed ?

__global__ void example_kernel() {
    enum class WarpRole {
        ....
    };
    const auto num_sms = static_cast<int>(gridDim.x);
    const auto sm_id = static_cast<int>(blockIdx.x);
    const auto thread_id = static_cast<int>(threadIdx.x), warp_id = thread_id / 32, lane_id = get_lane_id();
    const bool is_forwarder = sm_id % 2 == 0;
    const auto role_meta = [=]() -> WarpRole {
        ....
    }();
    auto warp_role = role_meta;
    if (warp_role == WarpRole::kRDMAAndNVLForwarder) {
        printf("Warp %d on SM %d Forwarder\n", warp_id, sm_id);
        trap();
    } else if (warp_role == WarpRole::kForwarderCoordinator) {
        printf("Warp %d on SM %d is a Forwarder_Coordinator\n", warp_id, sm_id);
        trap();
    } else if (warp_role == WarpRole::kRDMASender) {
        printf("Warp %d on SM %d is a RDMA_Sender\n", warp_id, sm_id);
        trap();
    } else if (warp_role == WarpRole::kRDMASenderCoordinator) {
        printf("Warp %d on SM %d is a RDMA_Sender_Coordinator\n", warp_id, sm_id);
        trap();
    } else {
        printf("Warp %d on SM %d is NVL_Receivers\n", warp_id, sm_id);
        trap();
    }
}
const int block_size = 512;
const int grid_size = 24;
example_kernel<<<grid_size, block_size>>>();
cudaDeviceSynchronize();

Originally created by @huakainomore on GitHub (Nov 5, 2025). Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/480 ### 1. Description I encountered some errors just like this. Rdma_rank 1 have this error, the sender coordinator and receiver both have timeout error. 1. For NVL receicer, `head == tail`, it seem to be related to Forwarder. 2. For Sender coordinator, It depends on Sender to update rdma_send_channel_tail , so it seem to be related to Sender. but both the forwarder and Sender does not, we can't know the forwarder or sender has error or not ? I think it's related to trap(). After trap being called by on warp, The execution of the kernel is aborted, if other warp's print buffer can't be flushed to output, it may be lost. ``` DeepEP dispatch NVL receiver timeout, channel: 1, RDMA: 1, nvl: 1, src NVL: 2, head: 75, tail: 75 DeepEP RDMA sender coordinator timeout, channel: 1, IB: 1, nvl 2, dst IB: 0, tail: 128, remaining: 21 DeepEP RDMA sender coordinator timeout, channel: 1, IB: 1, nvl 2, dst IB: 1, tail: 128, remaining: 27 DeepEP RDMA sender coordinator timeout, channel: 1, IB: 1, nvl 2, dst IB: 2, tail: 128, remaining: 30 DeepEP RDMA sender coordinator timeout, channel: 1, IB: 1, nvl 2, dst IB: 3, tail: 112, remaining: 25 DeepEP dispatch NVL receiver timeout, channel: 1, RDMA: 1, nvl: 7, src NVL: 2, head: 85, tail: 85 DeepEP dispatch NVL receiver timeout, channel: 1, RDMA: 1, nvl: 6, src NVL: 2, head: 83, tail: 83 DeepEP dispatch NVL receiver timeout, channel: 1, RDMA: 1, nvl: 5, src NVL: 2, head: 88, tail: 88 DeepEP dispatch NVL receiver timeout, channel: 1, RDMA: 1, nvl: 0, src NVL: 2, head: 76, tail: 76 DeepEP dispatch NVL receiver timeout, channel: 1, RDMA: 1, nvl: 4, src NVL: 2, head: 86, tail: 86 DeepEP dispatch NVL receiver timeout, channel: 1, RDMA: 1, nvl: 3, src NVL: 2, head: 84, tail: 84 ``` ### 2. Verification Just to verify the phenomenon exist, the program can't output the all warp role, we can only see few warp output in same SM. I add some sleep before trap, it seem all warp have output, But it may not be guaranteed. So is there a way to ensure that all warp outputs are guaranteed ? ``` __global__ void example_kernel() { enum class WarpRole { .... }; const auto num_sms = static_cast<int>(gridDim.x); const auto sm_id = static_cast<int>(blockIdx.x); const auto thread_id = static_cast<int>(threadIdx.x), warp_id = thread_id / 32, lane_id = get_lane_id(); const bool is_forwarder = sm_id % 2 == 0; const auto role_meta = [=]() -> WarpRole { .... }(); auto warp_role = role_meta; if (warp_role == WarpRole::kRDMAAndNVLForwarder) { printf("Warp %d on SM %d Forwarder\n", warp_id, sm_id); trap(); } else if (warp_role == WarpRole::kForwarderCoordinator) { printf("Warp %d on SM %d is a Forwarder_Coordinator\n", warp_id, sm_id); trap(); } else if (warp_role == WarpRole::kRDMASender) { printf("Warp %d on SM %d is a RDMA_Sender\n", warp_id, sm_id); trap(); } else if (warp_role == WarpRole::kRDMASenderCoordinator) { printf("Warp %d on SM %d is a RDMA_Sender_Coordinator\n", warp_id, sm_id); trap(); } else { printf("Warp %d on SM %d is NVL_Receivers\n", warp_id, sm_id); trap(); } } const int block_size = 512; const int grid_size = 24; example_kernel<<<grid_size, block_size>>>(); cudaDeviceSynchronize(); ```

bsakaguchi commented

2026-10-26 09:58:22 +00:00

Author

Owner

@gangxie112 commented on GitHub (Nov 6, 2025):

@sphish Could help have a look at this？

@gangxie112 commented on GitHub (Nov 6, 2025): @sphish Could help have a look at this？

bsakaguchi commented

2026-10-26 09:58:23 +00:00

Author

Owner

@sphish commented on GitHub (Nov 6, 2025):

You can try replacing trap() with while(true) — this might help you output more logs.

@sphish commented on GitHub (Nov 6, 2025): You can try replacing `trap()` with `while(true)` — this might help you output more logs.

bsakaguchi commented

2026-10-26 09:58:23 +00:00

Author

Owner

@huakainomore commented on GitHub (Nov 6, 2025):

You can try replacing trap() with while(true) — this might help you output more logs.

Thanks for replay very much ! Will there be some fix to ensure that all warp outputs are guaranteed ?

@huakainomore commented on GitHub (Nov 6, 2025): > You can try replacing `trap()` with `while(true)` — this might help you output more logs. Thanks for replay very much ! Will there be some fix to ensure that all warp outputs are guaranteed ?

bsakaguchi commented

2026-10-26 09:58:23 +00:00

Author

Owner

@gangxie112 commented on GitHub (Nov 6, 2025):

You can try replacing trap() with while(true) — this might help you output more logs.

I don't think it could work. Because this is running on our production env. When inifite loop，we have to wait for some other timeout triggered, like pytorch nccl-watchdog, which usually takes more time.

@gangxie112 commented on GitHub (Nov 6, 2025): > You can try replacing `trap()` with `while(true)` — this might help you output more logs. I don't think it could work. Because this is running on our production env. When inifite loop，we have to wait for some other timeout triggered, like pytorch nccl-watchdog, which usually takes more time.

bsakaguchi commented

2026-10-26 09:58:23 +00:00

Author

Owner

@sphish commented on GitHub (Nov 6, 2025):

I’m not sure if there’s a better solution — in fact, this issue often troubles me as well.
Debugging in a production environment is very challenging; it would be best if you could reproduce this problem using a unit test.

@sphish commented on GitHub (Nov 6, 2025): I’m not sure if there’s a better solution — in fact, this issue often troubles me as well. Debugging in a production environment is very challenging; it would be best if you could reproduce this problem using a unit test.

bsakaguchi commented

2026-10-26 09:58:24 +00:00

Author

Owner

@huakainomore commented on GitHub (Nov 7, 2025):

I’m not sure if there’s a better solution — in fact, this issue often troubles me as well. Debugging in a production environment is very challenging; it would be best if you could reproduce this problem using a unit test.

I add some wait before trap()

__device__ __forceinline__ void trap() {
    auto end_time = clock64() + NUM_TIMEOUT_CYCLES;
    while (clock64() < end_time) {
        __nanosleep(NUM_WAIT_NANOSECONDS);
    }
    asm("trap;");
}

add break for lane 0 and channel 5 kForwarderCoordinator to make timeout

       // kForwarderCoordinator
        while (true) {
            // Find minimum head
            int min_head = std::numeric_limits<int>::max();
            #pragma unroll
            for (int i = 0; i < NUM_MAX_NVL_PEERS; ++ i) if (not forward_channel_retired[i])
                min_head = min(min_head, forward_channel_head[i][target_rdma]);
            if (__all_sync(0xffffffff, min_head == std::numeric_limits<int>::max()))
                break;

            if (lane_id == 0 && channel_id == 5) {
                break;
            }

the origin result the rdma_rank 0 only have RDMA sender coordinator timeout, after add wait before trap(), the all WarpRole will have all timeout log because the double NUM_TIMEOUT_CYCLES wait time.

I think this is a fix for better locate problem, the double NUM_TIMEOUT_CYCLES wait time may can help other WarpRole output.
@sphish

@huakainomore commented on GitHub (Nov 7, 2025): > I’m not sure if there’s a better solution — in fact, this issue often troubles me as well. Debugging in a production environment is very challenging; it would be best if you could reproduce this problem using a unit test. I add some wait before trap() ``` __device__ __forceinline__ void trap() { auto end_time = clock64() + NUM_TIMEOUT_CYCLES; while (clock64() < end_time) { __nanosleep(NUM_WAIT_NANOSECONDS); } asm("trap;"); } ``` add break for lane 0 and channel 5 kForwarderCoordinator to make timeout ``` // kForwarderCoordinator while (true) { // Find minimum head int min_head = std::numeric_limits<int>::max(); #pragma unroll for (int i = 0; i < NUM_MAX_NVL_PEERS; ++ i) if (not forward_channel_retired[i]) min_head = min(min_head, forward_channel_head[i][target_rdma]); if (__all_sync(0xffffffff, min_head == std::numeric_limits<int>::max())) break; if (lane_id == 0 && channel_id == 5) { break; } ``` the origin result the rdma_rank 0 only have RDMA sender coordinator timeout, after add wait before trap(), the all WarpRole will have all timeout log because the double NUM_TIMEOUT_CYCLES wait time. I think this is a fix for better locate problem, the double NUM_TIMEOUT_CYCLES wait time may can help other WarpRole output. @sphish

bsakaguchi commented

2026-10-26 09:58:24 +00:00

Author

Owner

@1998zxn commented on GitHub (Nov 10, 2025):

I noticed that the CPU timeout and GPU timeout cycles were set in the configuration file. After We changed trap () to while(true), We increased the CPU timeout, which also helped us print all timeout logs before exiting the program.

@1998zxn commented on GitHub (Nov 10, 2025): I noticed that the CPU timeout and GPU timeout cycles were set in the configuration file. After We changed trap () to while(true), We increased the CPU timeout, which also helped us print all timeout logs before exiting the program.

bsakaguchi commented

2026-10-26 09:58:24 +00:00

Author

Owner

@huakainomore commented on GitHub (Nov 11, 2025):

I noticed that the CPU timeout and GPU timeout cycles were set in the configuration file. After We changed trap () to while(true), We increased the CPU timeout, which also helped us print all timeout logs before exiting the program.

Hi, thanks for reply! i have a question for it, if use cached_notify in internode_dispatch, it will not check CPU timeout, after replace trap() to while(true), the program wil nerver exiting?

@huakainomore commented on GitHub (Nov 11, 2025): > I noticed that the CPU timeout and GPU timeout cycles were set in the configuration file. After We changed trap () to while(true), We increased the CPU timeout, which also helped us print all timeout logs before exiting the program. Hi, thanks for reply! i have a question for it, if use cached_notify in internode_dispatch, it will not check CPU timeout, after replace trap() to while(true), the program wil nerver exiting?

Rows
Columns

[GH-ISSUE #480] [bug]Missing some critical error log when timeout #309

1. Description

2. Verification