[GH-ISSUE #549] [Question] tma_store_wait waits for multiple consecutive TMA tasks issued by the current thread #341

New issue

Open

opened 2026-10-26 09:58:50 +00:00 by bsakaguchi · 0 comments

bsakaguchi commented

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

Owner

Originally created by @cyhdmjzzy on GitHub (Dec 26, 2025).
Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/549

In the receiver side of combine in intranode.cu when using TMA, it should be:

if (lane_id < kNumStages)

But currently it is:

if (elect_one_sync())

This is because: Each invocation of tma_store_1d issues a TMA bulk_group, and multiple TMA bulk_group issued by the same thread execute serially (see NVIDIA PTX documentation: 9.7.9.27.2.2. Data Movement and Conversion Instructions: cp.async.bulk.wait_group).
The specific changes are located in PR 527: deepseek-ai/DeepEP

Originally created by @cyhdmjzzy on GitHub (Dec 26, 2025). Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/549 In the receiver side of `combine` in `intranode.cu` when using TMA, it should be: ```cpp if (lane_id < kNumStages) ``` But currently it is: ```cpp if (elect_one_sync()) ``` This is because: Each invocation of `tma_store_1d` issues a TMA `bulk_group`, and multiple TMA `bulk_group` issued by the same thread execute serially (see [NVIDIA PTX documentation: 9.7.9.27.2.2. Data Movement and Conversion Instructions: cp.async.bulk.wait_group](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-wait-group)). The specific changes are located in [PR 527: deepseek-ai/DeepEP](https://github.com/deepseek-ai/DeepEP/pull/527/files)

bsakaguchi referenced this issue

2026-10-27 21:03:41 +00:00

[PR #341] [MERGED] Support nvfp4 low latency mode dispatch #526

bsakaguchi referenced this issue

2026-10-27 21:03:58 +00:00

[PR #405] One branch that contains several optimizations and features #551