deepseek-ai / deepseek-ai/DeepEP
[Question] tma_store_wait waits for multiple consecutive TMA tasks issued by the current thread
- 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