deepseek-ai/DeepGEMM
1
0
Fork 0
mirror of https://github.com/deepseek-ai/DeepGEMM.git synced 2026-11-05 00:47:26 +00:00
Code Issues 162 Projects Packages Wiki Activity

[PR #375] Fix non-deterministic results in k_grouped_fp8_gemm_nt_contiguous #397

New issue
Open
opened 2026-10-26 21:17:13 +00:00 by bsakaguchi · 0 comments
bsakaguchi commented 2026-10-26 21:17:13 +00:00
Owner
Copy link

📋 Pull Request Information

Original PR: https://github.com/deepseek-ai/DeepGEMM/pull/375
Author: @Functionhx
Created: 7/9/2026
Status: 🔄 Open

Base: main ← Head: fix/fp8-gemm-nondeterministic

📝 Commits (1)

  • 4e4e540 Fix non-deterministic numerical results in k_grouped_fp8_gemm_nt_contiguous

📊 Changes

1 file changed (+10 additions, -7 deletions)

View changed files

📝 deep_gemm/include/deep_gemm/impls/sm90_fp8_gemm_1d1d.cuh (+10 -7)

📄 Description

Summary

Fixes #365: non-deterministic numerical results in FP8 grouped GEMM when one expert has exactly 128 tokens.

Root Cause

Race condition in sm90_fp8_gemm_1d1d.cuh: tensormap.replace operations and in-flight TMA data loads were in the same cp.async.bulk commit group. Per PTX ISA, operations within a commit group may complete in any order, allowing descriptor modifications to become visible before pending TMA reads complete.

Fix

Restructure group-switch into three phases:

  1. Drain: commit_group + wait_group all pending TMA data operations
  2. Modify + Fence: issue tensormap.replace, then commit_group + wait_group in a separate group
  3. Publish: copy updated descriptors to global memory, release_gpu + acquire_gpu

🔄 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/DeepGEMM/pull/375 **Author:** [@Functionhx](https://github.com/Functionhx) **Created:** 7/9/2026 **Status:** 🔄 Open **Base:** `main` ← **Head:** `fix/fp8-gemm-nondeterministic` --- ### 📝 Commits (1) - [`4e4e540`](https://github.com/deepseek-ai/DeepGEMM/commit/4e4e54036becc2a1ee9b5470015566673045acfd) Fix non-deterministic numerical results in k_grouped_fp8_gemm_nt_contiguous ### 📊 Changes **1 file changed** (+10 additions, -7 deletions) <details> <summary>View changed files</summary> 📝 `deep_gemm/include/deep_gemm/impls/sm90_fp8_gemm_1d1d.cuh` (+10 -7) </details> ### 📄 Description ## Summary Fixes #365: non-deterministic numerical results in FP8 grouped GEMM when one expert has exactly 128 tokens. ## Root Cause Race condition in sm90_fp8_gemm_1d1d.cuh: tensormap.replace operations and in-flight TMA data loads were in the same cp.async.bulk commit group. Per PTX ISA, operations within a commit group may complete in any order, allowing descriptor modifications to become visible before pending TMA reads complete. ## Fix Restructure group-switch into three phases: 1. Drain: commit_group + wait_group all pending TMA data operations 2. Modify + Fence: issue tensormap.replace, then commit_group + wait_group in a separate group 3. Publish: copy updated descriptors to global memory, release_gpu + acquire_gpu --- <sub>🔄 This issue represents a GitHub Pull Request. It cannot be merged through Gitea due to API limitations.</sub>
Sign in to join this conversation.
No Branch/Tag specified
Branches Tags
main
public-sync-20261014
0930-release
nv_dev
fix-mega-moe-fence
public-release-2609
public-release-2607
public-release-2606
mega-update
public-release-260416
ko3n1g/ci/build-wheels
nv_dev_f8e8fb5
nv_dev_950d31a
nv_dev_67fc648
nv_dev_c491439
nv_dev_4ff3f54
v2.1.1.post3
v2.1.1.post2
v2.1.1.post1
v2.1.1
v2.1.0
Labels
Clear labels   
pull-request
Mirrored from GitHub Pull Request

No labels
pull-request
Milestone
Clear milestone
No items
No milestone
Projects
Clear projects
No items
No project
Assignees
Clear assignees
No assignees
1 participant
Notifications
Due date
The due date is invalid or out of range. Please use the format "yyyy-mm-dd".

No due date set.

Dependencies

No dependencies set.

Reference
deepseek-ai/DeepGEMM#397
No description provided.