[GH-ISSUE #481] Does DeepEP support sm86?(4A6000 without NVL) #310

New issue

Open

opened 2026-10-26 09:58:25 +00:00 by bsakaguchi · 9 comments

bsakaguchi commented 2026-10-26 09:58:25 +00:00

Owner

Copy link

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

Hello DeepEP Team,

System Environment
GPU: 4x NVIDIA A6000 (Ampere Architecture, Compute Capability sm_86)

Interconnect: GPUs are connected via PCIe (no NVLink).

Problem Description
export TORCH_CUDA_ARCH_LIST="8.6"
export DISABLE_SM90_FEATURES="1"
export NVSHMEM_DIR=$(...)
export CPLUS_INCLUDE_PATH=$CONDA_PREFIX/include
export LDFLAGS="-L/usr/local/cuda/lib64/stubs"

python setup.py install

** Error **

adding license file 'LICENSE'
writing manifest file 'deep_ep.egg-info/SOURCES.txt'
Copying deep_ep.egg-info to /home/zyhuang/miniconda3/envs/deepep_env/lib/python3.10/site-packages/deep_ep-1.2.1+9249c25-py3.10.egg-info
running install_scripts
(deepep_env) ada6000-02% python
Python 3.10.19 | packaged by conda-forge | (main, Oct 22 2025, 22:29:10) [GCC 14.3.0] on linux
Type "help", "copyright", "credits" or "license" for more information.
>>> import deep_ep
Traceback (most recent call last):
  File "<stdin>", line 1, in <module>
  File "/home/zyhuang/temp_can/distributed/DeepEP/deep_ep/__init__.py", line 3, in <module>
    from .utils import EventOverlap
  File "/home/zyhuang/temp_can/distributed/DeepEP/deep_ep/utils.py", line 7, in <module>
    from deep_ep_cpp import EventHandle
ImportError: /home/zyhuang/miniconda3/envs/deepep_env/lib/python3.10/site-packages/deep_ep_cpp.cpython-310-x86_64-linux-gnu.so: undefined symbol: _ZN7deep_ep12internode_ll17query_mask_bufferEPiiS1_P11CUstream_st

====> I can install, but can not import

Originally created by @ziyuhuang123 on GitHub (Nov 5, 2025). Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/481 Hello DeepEP Team, **System Environment** GPU: 4x NVIDIA A6000 (Ampere Architecture, Compute Capability sm_86) Interconnect: GPUs are connected via PCIe (no NVLink). **Problem Description** export TORCH_CUDA_ARCH_LIST="8.6" export DISABLE_SM90_FEATURES="1" export NVSHMEM_DIR=$(...) export CPLUS_INCLUDE_PATH=$CONDA_PREFIX/include export LDFLAGS="-L/usr/local/cuda/lib64/stubs" python setup.py install ** Error ** ``` adding license file 'LICENSE' writing manifest file 'deep_ep.egg-info/SOURCES.txt' Copying deep_ep.egg-info to /home/zyhuang/miniconda3/envs/deepep_env/lib/python3.10/site-packages/deep_ep-1.2.1+9249c25-py3.10.egg-info running install_scripts (deepep_env) ada6000-02% python Python 3.10.19 | packaged by conda-forge | (main, Oct 22 2025, 22:29:10) [GCC 14.3.0] on linux Type "help", "copyright", "credits" or "license" for more information. >>> import deep_ep Traceback (most recent call last): File "<stdin>", line 1, in <module> File "/home/zyhuang/temp_can/distributed/DeepEP/deep_ep/__init__.py", line 3, in <module> from .utils import EventOverlap File "/home/zyhuang/temp_can/distributed/DeepEP/deep_ep/utils.py", line 7, in <module> from deep_ep_cpp import EventHandle ImportError: /home/zyhuang/miniconda3/envs/deepep_env/lib/python3.10/site-packages/deep_ep_cpp.cpython-310-x86_64-linux-gnu.so: undefined symbol: _ZN7deep_ep12internode_ll17query_mask_bufferEPiiS1_P11CUstream_st ``` ====> I can install, but can not import

bsakaguchi commented 2026-10-26 09:58:26 +00:00

Author

Owner

Copy link

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

Normal kernels do not support; low‑latency kernels might support your system.

The issue you’re encountering is that you need to specify the correct TORCH_CUDA_ARCH_LIST during installation.

<!-- gh-comment-id:3494924092 --> @sphish commented on GitHub (Nov 6, 2025): Normal kernels do not support; low‑latency kernels might support your system. The issue you’re encountering is that you need to specify the correct `TORCH_CUDA_ARCH_LIST` during installation.

bsakaguchi commented 2026-10-26 09:58:26 +00:00

Author

Owner

Copy link

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

same error ImportError: /root/miniconda3/envs/py310/lib/python3.10/site-packages/deep_ep_cpp.cpython-310-x86_64-linux-gnu.so: undefined symbol: _ZN7deep_ep12internode_ll17query_mask_bufferEPiiS1_P11CUstream_st
I install it with flag TORCH_CUDA_ARCH_LIST=8.0 DISABLE_SM90_FEATURES=1 on L20*8 server

<!-- gh-comment-id:3500418839 --> @chenhongyu2048 commented on GitHub (Nov 7, 2025): same error `ImportError: /root/miniconda3/envs/py310/lib/python3.10/site-packages/deep_ep_cpp.cpython-310-x86_64-linux-gnu.so: undefined symbol: _ZN7deep_ep12internode_ll17query_mask_bufferEPiiS1_P11CUstream_st` I install it with flag `TORCH_CUDA_ARCH_LIST=8.0 DISABLE_SM90_FEATURES=1` on L20*8 server

bsakaguchi commented 2026-10-26 09:58:27 +00:00

Author

Owner

Copy link

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

seems this is because in setup.py, deep_ep disable internode and low-latency kernels, but in some others code these kernels are used

<!-- gh-comment-id:3500529234 --> @chenhongyu2048 commented on GitHub (Nov 7, 2025): seems this is because in `setup.py`, deep_ep disable internode and low-latency kernels, but in some others code these kernels are used

bsakaguchi commented 2026-10-26 09:58:27 +00:00

Author

Owner

Copy link

@LinglingGreat commented on GitHub (Nov 12, 2025):

same error xxx/lib/python3.10/site-packages/deep_ep_cpp.cpython-310-x86_64-linux-gnu.so: undefined symbol: _ZN7deep_ep12internode_ll17query_mask_bufferEPiiS1_P11CUstream_st

Use TORCH_CUDA_ARCH_LIST=8.0 DISABLE_SM90_FEATURES=1 on A100

<!-- gh-comment-id:3521350673 --> @LinglingGreat commented on GitHub (Nov 12, 2025): same error `xxx/lib/python3.10/site-packages/deep_ep_cpp.cpython-310-x86_64-linux-gnu.so: undefined symbol: _ZN7deep_ep12internode_ll17query_mask_bufferEPiiS1_P11CUstream_st` Use `TORCH_CUDA_ARCH_LIST=8.0 DISABLE_SM90_FEATURES=1` on A100

bsakaguchi commented 2026-10-26 09:58:27 +00:00

Author

Owner

Copy link

@zhuoshanzhou commented on GitHub (Nov 12, 2025):

I also encountered this issue, and it can be resolved by modifying the code from github.com/deepseek-ai/DeepEP@92fe2deaec/csrc/deep_ep.cpp (L1825) to github.com/deepseek-ai/DeepEP@92fe2deaec/csrc/deep_ep.cpp (L1842) with

#ifndef DISABLE_NVSHMEM

#else
    EP_HOST_ASSERT(false and "NVSHMEM is disabled during compilation");
#endif

In addition, another issue I met is that cudaIpcMemLazyEnablePeerAccess does not always work, so I explicitly enabled P2P in the Python test function.

<!-- gh-comment-id:3521776264 --> @zhuoshanzhou commented on GitHub (Nov 12, 2025): I also encountered this issue, and it can be resolved by modifying the code from https://github.com/deepseek-ai/DeepEP/blob/92fe2deaec24bc92ebd9de276daa6ca9ed602ed4/csrc/deep_ep.cpp#L1825 to https://github.com/deepseek-ai/DeepEP/blob/92fe2deaec24bc92ebd9de276daa6ca9ed602ed4/csrc/deep_ep.cpp#L1842 with ```cpp #ifndef DISABLE_NVSHMEM #else EP_HOST_ASSERT(false and "NVSHMEM is disabled during compilation"); #endif ``` In addition, another issue I met is that cudaIpcMemLazyEnablePeerAccess does not always work, so I explicitly enabled P2P in the Python test function.

bsakaguchi commented 2026-10-26 09:58:27 +00:00

Author

Owner

Copy link

@BarrinXu commented on GitHub (Nov 14, 2025):

Hi @zhuoshanzhou , I also encountered the same issue. Could you please share more details on how you resolve this problem? I cannot understand the two lines code you mentioned. Should we replace these two functions' inside code with

#ifndef DISABLE_NVSHMEM

#else
    EP_HOST_ASSERT(false and "NVSHMEM is disabled during compilation");
#endif

Many thanks in advance.

upd: I have resolved this problem by replacing with the code you provide, thanks.

<!-- gh-comment-id:3531922515 --> @BarrinXu commented on GitHub (Nov 14, 2025): Hi @zhuoshanzhou , I also encountered the same issue. Could you please share more details on how you resolve this problem? I cannot understand the two lines code you mentioned. Should we replace these two functions' inside code with ``` #ifndef DISABLE_NVSHMEM #else EP_HOST_ASSERT(false and "NVSHMEM is disabled during compilation"); #endif ``` Many thanks in advance. upd: I have resolved this problem by replacing with the code you provide, thanks.

bsakaguchi commented 2026-10-26 09:58:27 +00:00

Author

Owner

Copy link

@bleedingfight commented on GitHub (Nov 28, 2025):

@BarrinXu how to replace it？

void Buffer::low_latency_query_mask_buffer(const torch::Tensor& mask_status) {
#ifndef DISABLE_NVSHMEM
    EP_HOST_ASSERT(mask_buffer_ptr != nullptr and "Shrink mode must be enabled");
    EP_HOST_ASSERT(mask_status.numel() == num_ranks && mask_status.scalar_type() == torch::kInt32);
    internode_ll::query_mask_buffer(
        mask_buffer_ptr, num_ranks, reinterpret_cast<int*>(mask_status.data_ptr()), at::cuda::getCurrentCUDAStream());

#else
    EP_HOST_ASSERT(false and "NVSHMEM is disabled during compilation");
#endif
}

？

<!-- gh-comment-id:3588384717 --> @bleedingfight commented on GitHub (Nov 28, 2025): @BarrinXu how to replace it？ ```cpp void Buffer::low_latency_query_mask_buffer(const torch::Tensor& mask_status) { #ifndef DISABLE_NVSHMEM EP_HOST_ASSERT(mask_buffer_ptr != nullptr and "Shrink mode must be enabled"); EP_HOST_ASSERT(mask_status.numel() == num_ranks && mask_status.scalar_type() == torch::kInt32); internode_ll::query_mask_buffer( mask_buffer_ptr, num_ranks, reinterpret_cast<int*>(mask_status.data_ptr()), at::cuda::getCurrentCUDAStream()); #else EP_HOST_ASSERT(false and "NVSHMEM is disabled during compilation"); #endif } ``` ？

bsakaguchi commented 2026-10-26 09:58:27 +00:00

Author

Owner

Copy link

@BarrinXu commented on GitHub (Nov 28, 2025):

@BarrinXu how to replace it？

void Buffer::low_latency_query_mask_buffer(const torch::Tensor& mask_status) {
#ifndef DISABLE_NVSHMEM
EP_HOST_ASSERT(mask_buffer_ptr != nullptr and "Shrink mode must be enabled");
EP_HOST_ASSERT(mask_status.numel() == num_ranks && mask_status.scalar_type() == torch::kInt32);
internode_ll::query_mask_buffer(
mask_buffer_ptr, num_ranks, reinterpret_cast<int*>(mask_status.data_ptr()), at::cuda::getCurrentCUDAStream());

#else
EP_HOST_ASSERT(false and "NVSHMEM is disabled during compilation");
#endif
}
？

@bleedingfight you can try this

void Buffer::low_latency_update_mask_buffer(int rank_to_mask, bool mask) {
    #ifndef DISABLE_NVSHMEM

    #else
        EP_HOST_ASSERT(false and "NVSHMEM is disabled during compilation");
    #endif
}

void Buffer::low_latency_query_mask_buffer(const torch::Tensor& mask_status) {
    #ifndef DISABLE_NVSHMEM

    #else
        EP_HOST_ASSERT(false and "NVSHMEM is disabled during compilation");
    #endif
}

void Buffer::low_latency_clean_mask_buffer() {
    #ifndef DISABLE_NVSHMEM

    #else
        EP_HOST_ASSERT(false and "NVSHMEM is disabled during compilation");
    #endif
}

<!-- gh-comment-id:3588458748 --> @BarrinXu commented on GitHub (Nov 28, 2025): > [@BarrinXu](https://github.com/BarrinXu) how to replace it？ > > void Buffer::low_latency_query_mask_buffer(const torch::Tensor& mask_status) { > #ifndef DISABLE_NVSHMEM > EP_HOST_ASSERT(mask_buffer_ptr != nullptr and "Shrink mode must be enabled"); > EP_HOST_ASSERT(mask_status.numel() == num_ranks && mask_status.scalar_type() == torch::kInt32); > internode_ll::query_mask_buffer( > mask_buffer_ptr, num_ranks, reinterpret_cast<int*>(mask_status.data_ptr()), at::cuda::getCurrentCUDAStream()); > > #else > EP_HOST_ASSERT(false and "NVSHMEM is disabled during compilation"); > #endif > } > ？ @bleedingfight you can try this ``` void Buffer::low_latency_update_mask_buffer(int rank_to_mask, bool mask) { #ifndef DISABLE_NVSHMEM #else EP_HOST_ASSERT(false and "NVSHMEM is disabled during compilation"); #endif } void Buffer::low_latency_query_mask_buffer(const torch::Tensor& mask_status) { #ifndef DISABLE_NVSHMEM #else EP_HOST_ASSERT(false and "NVSHMEM is disabled during compilation"); #endif } void Buffer::low_latency_clean_mask_buffer() { #ifndef DISABLE_NVSHMEM #else EP_HOST_ASSERT(false and "NVSHMEM is disabled during compilation"); #endif } ```

bsakaguchi commented 2026-10-26 09:58:27 +00:00

Author

Owner

Copy link

@bleedingfight commented on GitHub (Dec 1, 2025):

@BarrinXu thx for your replay，IAfter using DISABLE-NVSHMEM, the problem was solved！

<!-- gh-comment-id:3594193616 --> @bleedingfight commented on GitHub (Dec 1, 2025): @BarrinXu thx for your replay，IAfter using DISABLE-NVSHMEM, the problem was solved！

bsakaguchi referenced this issue 2026-10-27 21:03:31 +00:00

[PR #310] [MERGED] Fix for data error and kernel hung because of inflight rdma channel head update #515

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

Subscribe

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#310

No description provided.