Skip to content

[BugFix][sm90] tf32 gemm race condition - #448

Open
leevan wants to merge 1 commit into
deepseek-ai:mainfrom
leevan:fix/sm90-hc-prenorm-wgmma-race
Open

leevan wants to merge 1 commit into
deepseek-ai:mainfrom
leevan:fix/sm90-hc-prenorm-wgmma-race

Conversation

@leevan

@leevan leevan commented Sep 15, 2026

Copy link
Copy Markdown

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.

Fix a WGMMA register reuse race in tf32_hc_prenorm_gemm and add a regression test.
}
ptx::warpgroup_commit_batch();
// Keep register A operands live until the asynchronous WGMMA completes.
ptx::warpgroup_wait<0>();

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🔵 suggestion: Placing ptx::warpgroup_wait<0>() after warpgroup_commit_batch() is the correct minimal fix: it keeps the single-buffered register operands a[...] live until the async WGMMA has consumed them. Consequence is that every K stage now fully drains its WGMMA before the next A load, so WGMMA and the next stage's A-load/reduction cannot overlap. This is not a regression versus the buggy code (which already waited each iteration) and correctness must come first, but it is worth a short follow-up note: if the kernel is compute bound, double-buffering a and using warpgroup_wait<1>() would restore overlap while still protecting the register source.

🤖 v4


ptx::warpgroup_wait<0>();
if (s > 0)
empty_barriers[(s - 1) % kNumStages]->arrive();

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🔵 suggestion: With the wait moved to the end of the loop body, empty_barriers[(s - 1) % kNumStages]->arrive() no longer has an explicit preceding warpgroup_wait<0>() in this iteration; it relies on the wait performed at the end of the previous iteration (i.e. the stage s-1 WGMMA is already complete). That is correct today, but the dependency is now implicit and a future refactor could silently reintroduce the race. Please add a one-line comment next to the arrive stating that stage s-1's WGMMA has already been waited on, so it is safe to release smem_a/smem_b of that stage.

🤖 v4

# repeatability across repeated launches (a correct deterministic kernel
# must produce identical results every time).
previous_pdl = deep_gemm.get_pdl()
previous_tf32 = torch.backends.cuda.matmul.allow_tf32

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟡 warning: The test saves/restores torch.backends.cuda.matmul.allow_tf32, but test_hc_prenorm_gemm() (line 18) sets torch.backends.cudnn.allow_tf32 = True and never restores it. As a result the pre-existing test still leaks global cudnn TF32 state to the rest of the suite, and the new test does not fully isolate the backend flags either. Since the new test intentionally disables TF32 for the reference, please save/restore both torch.backends.cuda.matmul.allow_tf32 and torch.backends.cudnn.allow_tf32 (ideally in test_hc_prenorm_gemm() as well) so the two tests cannot affect each other.

🤖 v4

deep_gemm.set_pdl(enable_pdl)
stream = torch.cuda.Stream() if use_side_stream else torch.cuda.current_stream()
with torch.cuda.stream(stream):
generator = torch.Generator(device='cuda').manual_seed(123)

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🔵 suggestion: generator is recreated with manual_seed(123) inside all four nested loops, so each (pdl, side_stream) configuration runs on the exact same a/b values. That is fine for the repeatability assertion, but it means the sweep never exercises different data across configurations. Since the race is timing/data-order sensitive, deriving the seed from the loop variables (or using a different seed per configuration) would broaden coverage without hurting determinism.

🤖 v4

torch.testing.assert_close(sqr_sum, ref_s, rtol=1e-5, atol=2e-3)
if previous_d is not None:
# A correct deterministic kernel is bit-exact across launches.
torch.testing.assert_close(d, previous_d, rtol=0, atol=0)

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🔵 suggestion: assert_close(d, previous_d, rtol=0, atol=0) asserts bit-exact determinism across launches. The SM90 kernel is currently deterministic (outputs are partitioned per split with no atomics), so this holds, but the assertion is stronger than typical GEMM tests and would become a false failure if split-K reduction is ever changed to a non-deterministic scheme. Please add a brief rationale comment or a named helper (assert_bit_exact) so the intent and fragility are explicit.

🤖 v4

@ds-review-bot

Copy link
Copy Markdown
Collaborator

🤖 ds-review-bot Code Review

v6

⚠️ 未完成评审(budget_exceeded:模型额度已用尽)

v5

⚠️ 未完成评审(budget_exceeded:模型额度已用尽)

v4

The MR fixes a real SM90 register-source WGMMA race in sm90_tf32_hc_prenorm_gemm. Previously the kernel loaded the FP32 A fragment into single-buffered a[...] registers BEFORE waiting for the previous stage's asynchronous WGMMA to complete, so the in-flight WGMMA could read A registers that had already been overwritten. The fix moves ptx::warpgroup_wait&lt;0&gt;() from the top of the K loop (after the A load) to the end of the loop body, immediately after warpgroup_commit_batch(), guaranteeing the prior WGMMA is done before the next iteration reloads a[...]. This is correct and matches the wait discipline used by other SM90 kernels (e.g. sm90_bf16_gemm.cuh), and the empty_barriers[(s-1)]-&gt;arrive() ordering remains safe because the previous stage's WGMMA has now been drained. The new regression test is well targeted: it sweeps split-K, large K, PDL on/off and side streams, validates per-split results against a full-precision (TF32-disabled) reference, and asserts bit-exact repeatability across launches. I could not execute the test (no SM90 GPU in this environment), but the patch is logically sound. The A-operand wait at line 206 no longer has an explicit preceding wait because it now relies on the wait at the end of the previous iteration, which is correct but implicit. Remaining feedback is minor/non-blocking.

Files reviewed: 2
Issues found: 🟡 1 warning | 🔵 4 suggestion
Inline comments posted: 5

⚠️ Parse warning: [v6] budget_exceeded:模型额度已用尽; [v5] budget_exceeded:模型额度已用尽

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants