[PR #448] [BugFix][sm90] tf32 gemm race condition #453

Open
opened 2026-10-26 21:18:02 +00:00 by bsakaguchi · 0 comments
Owner

📋 Pull Request Information

Original PR: https://github.com/deepseek-ai/DeepGEMM/pull/448
Author: @leevan
Created: 9/15/2026
Status: 🔄 Open

Base: main ← Head: fix/sm90-hc-prenorm-wgmma-race


📝 Commits (1)

📊 Changes

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

View changed files

📝 deep_gemm/include/deep_gemm/impls/sm90_tf32_hc_prenorm_gemm.cuh (+2 -1)
📝 tests/test_hyperconnection.py (+61 -0)

📄 Description

Fix a WGMMA register reuse race in SM90 tf32_hc_prenorm_gemm by waiting for completion before reusing A registers.

Add a regression test covering split-K, PDL, and side streams, with reference and repeatability checks.

Validation: tests/test_hyperconnection.py passed on an SM90 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/448 **Author:** [@leevan](https://github.com/leevan) **Created:** 9/15/2026 **Status:** 🔄 Open **Base:** `main` ← **Head:** `fix/sm90-hc-prenorm-wgmma-race` --- ### 📝 Commits (1) - [`6bc1ce3`](https://github.com/deepseek-ai/DeepGEMM/commit/6bc1ce37cdae4b10b621eb0d69ac61185f1d60eb) bug fix ### 📊 Changes **2 files changed** (+63 additions, -1 deletions) <details> <summary>View changed files</summary> 📝 `deep_gemm/include/deep_gemm/impls/sm90_tf32_hc_prenorm_gemm.cuh` (+2 -1) 📝 `tests/test_hyperconnection.py` (+61 -0) </details> ### 📄 Description Fix a WGMMA register reuse race in SM90 tf32_hc_prenorm_gemm by waiting for completion before reusing A registers. Add a regression test covering split-K, PDL, and side streams, with reference and repeatability checks. Validation: tests/test_hyperconnection.py passed on an SM90 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 labels
pull-request
No milestone
No project
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#453
No description provided.