deepseek-ai / deepseek-ai/DeepEP

[Question] tma_store_wait waits for multiple consecutive TMA tasks issued by the current thread

Open
#549 0 comments 0 reactions 0 assignees View on GitHub
Dominant language
Cuda
Stars
10.1k
Forks
1.4k
Avg merge
4d 1h
Merged PRs (30d)
2

Description

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)

Contributor guide

No contributing guide indexed for this repository

Research direction

Start in intranode.cu at the receiver side of combine when TMA is enabled, then compare the current condition with the changes in PR 527 and review the referenced NVIDIA PTX wait_group documentation. Done means the receiver waits according to the intended number of TMA stages for bulk groups issued by the current thread.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
distributed-systems
Issue type
Bug
Difficulty
2/5
Estimated time
1-3 hours
Activity status
Stale
Clarity
Clearly specified
Newbie friendliness
48/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.