[PR #642] [MERGED] Add fence.proxy.async.shared::cta between mbarrier wait and TMA load. #683

New issue

Closed

opened 2026-10-27 21:05:14 +00:00 by bsakaguchi · 0 comments

bsakaguchi commented

2026-10-27 21:05:14 +00:00

Owner

📋 Pull Request Information

Original PR: https://github.com/deepseek-ai/DeepEP/pull/642
Author: @RayWang96
Created: 5/24/2026
Status: ✅ Merged
Merged: 5/26/2026
Merged by: @sphish

Base: main ← Head: fix_fence

📝 Commits (2)

1185547 Add fence.proxy.async.shared::cta between mbarrier wait and TMA load.
0cf14f2 Move fence before "mbarrier_arrive"

📊 Changes

2 files changed (+10 additions, -1 deletions)

View changed files

📝 csrc/kernels/legacy/internode_ll.cu (+1 -1)
📝 csrc/kernels/legacy/utils.cuh (+9 -0)

📄 Description

This fixes the issue mentioned in https://github.com/deepseek-ai/DeepEP/issues/621. Compared with the fix in https://github.com/deepseek-ai/DeepEP/pull/632, this version uses a CTA-scope fence, which resolves the problem while preserving performance.

The data race here comes from LDS and TMA loads, so we only need fence.proxy.async.shared::cta; to prevent the TMA load from being issued before LDS has completed.

Thanks to elvircrn's script, I've tried my fix on NVL72, the result is correct and the performance is as good as the original implementation.

_{🔄 This issue represents a GitHub Pull Request. It cannot be merged through Gitea due to API limitations.}

## 📋 Pull Request Information **Original PR:** https://github.com/deepseek-ai/DeepEP/pull/642 **Author:** [@RayWang96](https://github.com/RayWang96) **Created:** 5/24/2026 **Status:** ✅ Merged **Merged:** 5/26/2026 **Merged by:** [@sphish](https://github.com/sphish) **Base:** `main` ← **Head:** `fix_fence` --- ### 📝 Commits (2) - [`1185547`](https://github.com/deepseek-ai/DeepEP/commit/1185547946bd549d804f8329442785c1976a93c8) Add fence.proxy.async.shared::cta between mbarrier wait and TMA load. - [`0cf14f2`](https://github.com/deepseek-ai/DeepEP/commit/0cf14f2811c5d9e880b1f1850b1bf6ddd5769ac6) Move fence before "mbarrier_arrive" ### 📊 Changes **2 files changed** (+10 additions, -1 deletions) <details> <summary>View changed files</summary> 📝 `csrc/kernels/legacy/internode_ll.cu` (+1 -1) 📝 `csrc/kernels/legacy/utils.cuh` (+9 -0) </details> ### 📄 Description This fixes the issue mentioned in https://github.com/deepseek-ai/DeepEP/issues/621. Compared with the fix in https://github.com/deepseek-ai/DeepEP/pull/632, this version uses a CTA-scope fence, which resolves the problem while preserving performance. The data race here comes from LDS and TMA loads, so we only need fence.proxy.async.shared::cta; to prevent the TMA load from being issued before LDS has completed. Thanks to [elvircrn](https://github.com/elvircrn)'s script, I've tried my fix on NVL72, the result is correct and the performance is as good as the original implementation. --- <sub>🔄 This issue represents a GitHub Pull Request. It cannot be merged through Gitea due to API limitations.</sub>

bsakaguchi

2026-10-27 21:05:14 +00:00