[GH-ISSUE #199] crash when set NVSHMEM_IBGDA_ENABLE_MULTI_PORT to 1 in case of dual-port RNIC #163

Open
opened 2026-10-26 09:55:29 +00:00 by bsakaguchi · 1 comment
Owner

Originally created by @yanminjia on GitHub (Jun 10, 2025).
Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/199

When I ran test_internode.py in case of dual-port, the environment variable NVSHMEM_IBGDA_ENABLE_MULTI_PORT is set to 1. Unfortunately, DeepEP crashed when create rdma team by calling nvshmem_team_split_strided(...) in the following code snippet (runtime.cu:init(...)):

  if (low_latency_mode and num_ranks > NUM_MAX_NVL_PEERS) {
        EP_HOST_ASSERT(cpu_rdma_team == NVSHMEM_TEAM_INVALID);
        EP_HOST_ASSERT(num_ranks % NUM_MAX_NVL_PEERS == 0);
        EP_HOST_ASSERT(nvshmem_team_split_strided(NVSHMEM_TEAM_WORLD, rank % NUM_MAX_NVL_PEERS, NUM_MAX_NVL_PEERS,
                                                  num_ranks / NUM_MAX_NVL_PEERS, &cpu_rdma_team_config, 0, &cpu_rdma_team) == 0);
        EP_HOST_ASSERT(cpu_rdma_team != NVSHMEM_TEAM_INVALID);
    }

And print below error message:

[/workspace/nvshmem/src/host/team/team_internal.cpp:1421] cuda failed with an illegal memory access was encountered
[/workspace/nvshmem/src/host/team/team_internal.cpp:1421] cuda failed with an illegal memory access was encountered
[/workspace/nvshmem/src/host/team/team_internal.cpp:1421] cuda failed with an illegal memory access was encountered
[/workspace/nvshmem/src/host/team/team_internal.cpp:1421] cuda failed with an illegal memory access was encountered
[/workspace/nvshmem/src/host/team/team_internal.cpp:1421] cuda failed with an illegal memory access was encountered
[/workspace/nvshmem/src/host/team/team_internal.cpp:1421] cuda failed with an illegal memory access was encountered
[/workspace/nvshmem/src/host/team/team_internal.cpp:1421] cuda failed with an illegal memory access was encountered
[/workspace/nvshmem/src/host/team/team_internal.cpp:1421] cuda failed with an illegal memory access was encountered

Any suggestion should be highly appreciated. Many thanks.

Originally created by @yanminjia on GitHub (Jun 10, 2025). Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/199 When I ran test_internode.py in case of dual-port, the environment variable _**NVSHMEM_IBGDA_ENABLE_MULTI_PORT**_ is set to 1. Unfortunately, DeepEP crashed when create rdma team by calling nvshmem_team_split_strided(...) in the following code snippet (runtime.cu:init(...)): ``` if (low_latency_mode and num_ranks > NUM_MAX_NVL_PEERS) { EP_HOST_ASSERT(cpu_rdma_team == NVSHMEM_TEAM_INVALID); EP_HOST_ASSERT(num_ranks % NUM_MAX_NVL_PEERS == 0); EP_HOST_ASSERT(nvshmem_team_split_strided(NVSHMEM_TEAM_WORLD, rank % NUM_MAX_NVL_PEERS, NUM_MAX_NVL_PEERS, num_ranks / NUM_MAX_NVL_PEERS, &cpu_rdma_team_config, 0, &cpu_rdma_team) == 0); EP_HOST_ASSERT(cpu_rdma_team != NVSHMEM_TEAM_INVALID); } ``` And print below error message: > [/workspace/nvshmem/src/host/team/team_internal.cpp:1421] cuda failed with an illegal memory access was encountered > [/workspace/nvshmem/src/host/team/team_internal.cpp:1421] cuda failed with an illegal memory access was encountered > [/workspace/nvshmem/src/host/team/team_internal.cpp:1421] cuda failed with an illegal memory access was encountered > [/workspace/nvshmem/src/host/team/team_internal.cpp:1421] cuda failed with an illegal memory access was encountered > [/workspace/nvshmem/src/host/team/team_internal.cpp:1421] cuda failed with an illegal memory access was encountered > [/workspace/nvshmem/src/host/team/team_internal.cpp:1421] cuda failed with an illegal memory access was encountered > [/workspace/nvshmem/src/host/team/team_internal.cpp:1421] cuda failed with an illegal memory access was encountered > [/workspace/nvshmem/src/host/team/team_internal.cpp:1421] cuda failed with an illegal memory access was encountered Any suggestion should be highly appreciated. Many thanks.
Author
Owner

@sphish commented on GitHub (Jun 10, 2025):

We haven’t used this environment variable before and are not familiar with how it works. I think you should consult the NVSHMEM team for more information.

<!-- gh-comment-id:2958499609 --> @sphish commented on GitHub (Jun 10, 2025): We haven’t used this environment variable before and are not familiar with how it works. I think you should consult the NVSHMEM team for more information.
Sign in to join this conversation.
No labels
pull-request
No milestone
No project
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#163
No description provided.