deepseek-ai/DeepEP

Fork 0

mirror of https://github.com/deepseek-ai/DeepEP.git synced 2026-11-05 00:47:32 +00:00

[GH-ISSUE #18] Reasons for not calling NVSHMEM public APIs? #18

New issue

Closed

opened 2026-10-26 09:52:31 +00:00 by bsakaguchi · 5 comments

bsakaguchi commented

2026-10-26 09:52:31 +00:00

Owner

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

First of all, this is an amazing work! Thank you so much for writing fantastic research papers in detail and open-sourcing the high performance building blocks!

As I was learning through the code base, I found that you folks wrote the kernels using NVSHMEM's IBGDA internals (nvshmemi_ibgda_*). I'm very curious what are the advantages of these internals over the public APIs (like nvshmemx_putmem_nbi_warp).

Originally created by @abcdabcd987 on GitHub (Feb 25, 2025). Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/18 First of all, this is an amazing work! Thank you so much for writing fantastic research papers in detail and open-sourcing the high performance building blocks! As I was learning through the code base, I found that you folks wrote the kernels using NVSHMEM's IBGDA internals (`nvshmemi_ibgda_*`). I'm very curious what are the advantages of these internals over the public APIs (like `nvshmemx_putmem_nbi_warp`).

bsakaguchi closed this issue

2026-10-26 09:52:31 +00:00

bsakaguchi commented

2026-10-26 09:52:32 +00:00

Author

Owner

@LyricZhao commented on GitHub (Feb 26, 2025):

For microsecond-level performance, the public APIs are not very good. For example, to ensure correctness, the public ones do a lot of sync and atomic operations. But our impls make each QP for a expert and do a lot of other assumptions, which can make the code much easier.

Also, the public API is not convenient for AR enabled (maybe you should do a quiet operation, which is slow), so for things like this, we can implement and co-design by rewriting these functions.

There are other optimizations in the NVSHMEM patch, e.g. shuffling the QP indices, all these can improve performance.

@LyricZhao commented on GitHub (Feb 26, 2025): For microsecond-level performance, the public APIs are not very good. For example, to ensure correctness, the public ones do a lot of sync and atomic operations. But our impls make each QP for a expert and do a lot of other assumptions, which can make the code much easier. Also, the public API is not convenient for AR enabled (maybe you should do a `quiet` operation, which is slow), so for things like this, we can implement and co-design by rewriting these functions. There are other optimizations in the NVSHMEM patch, e.g. shuffling the QP indices, all these can improve performance.

bsakaguchi commented

2026-10-26 09:52:33 +00:00

Author

Owner

@sphish commented on GitHub (Feb 26, 2025):

We aim to achieve finer-grained control over communication operations, as the native NVSHMEM interface has several limitations. One such limitation is the lack of support for fine-grained order-preserving operations. For example, when you need to send a flag after transmitting a piece of data, you can only rely on heavy operations like nvshmem_quiet(). Another example is our requirement to ensure each expert process exclusively uses a dedicated QP (Queue Pair), which cannot be achieved through NVSHMEM's native interfaces.

@sphish commented on GitHub (Feb 26, 2025): We aim to achieve finer-grained control over communication operations, as the native NVSHMEM interface has several limitations. One such limitation is the lack of support for fine-grained order-preserving operations. For example, when you need to send a flag after transmitting a piece of data, you can only rely on heavy operations like nvshmem_quiet(). Another example is our requirement to ensure each expert process exclusively uses a dedicated QP (Queue Pair), which cannot be achieved through NVSHMEM's native interfaces.

bsakaguchi commented

2026-10-26 09:52:33 +00:00

Author

Owner

@abcdabcd987 commented on GitHub (Feb 26, 2025):

Thanks for explanation! Makes lots of sense.

ensure each expert process exclusively uses a dedicated QP (Queue Pair)

I'm curious about this requirement. I know in the V3 paper, you use 1 expert per GPU. Do you support putting multiple experts on one GPU (for the low latency kernel)? And about dedicated QP, is it because you'd like to make sure that the GPU always uses the local NIC on the same PCIe switch, or is there some other consideration (like flow control etc.)

@abcdabcd987 commented on GitHub (Feb 26, 2025): Thanks for explanation! Makes lots of sense. > ensure each expert process exclusively uses a dedicated QP (Queue Pair) I'm curious about this requirement. I know in the V3 paper, you use 1 expert per GPU. Do you support putting multiple experts on one GPU (for the low latency kernel)? And about dedicated QP, is it because you'd like to make sure that the GPU always uses the local NIC on the same PCIe switch, or is there some other consideration (like flow control etc.)

bsakaguchi commented

2026-10-26 09:52:33 +00:00

Author

Owner

@sphish commented on GitHub (Feb 26, 2025):

In low-latency mode:

Different experts have no data dependencies, so utilizing different QPs (Queue Pairs) can enhance parallelism.
Data from the same expert must maintain order consistency, therefore communications for a single expert need to be processed within one QP.

@sphish commented on GitHub (Feb 26, 2025): In low-latency mode: 1. Different experts have no data dependencies, so utilizing different QPs (Queue Pairs) can enhance parallelism. 2. Data from the same expert must maintain order consistency, therefore communications for a single expert need to be processed within one QP.

bsakaguchi commented

2026-10-26 09:52:33 +00:00

Author

Owner

@sphish commented on GitHub (Feb 26, 2025):

And yes, we support putting multiple experts on one GPU.

@sphish commented on GitHub (Feb 26, 2025): And yes, we support putting multiple experts on one GPU.

bsakaguchi referenced this issue

2026-10-26 09:53:39 +00:00

[GH-ISSUE #74] test_internode.py hangs at internode.cu:dispatch(...) #68

bsakaguchi referenced this issue

2026-10-26 09:56:29 +00:00