deepseek-ai / deepseek-ai/DeepEP
[bug]Missing some critical error log when timeout
- Dominant language
- Cuda
- Stars
- 10.1k
- Forks
- 1.4k
- Avg merge
- 4d 1h
- Merged PRs (30d)
- 2
Description
### 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(gridDim.x);
const auto sm_id = static_cast(blockIdx.x);
const auto thread_id = static_cast(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<<>>();
cudaDeviceSynchronize();
```
Contributor guide
No contributing guide indexed for this repository
Assessment
This issue has not been assessed yet.