Skip to content

fix MMA Pipeline theoretical race#190

Open
Topologized wants to merge 1 commit intoHazyResearch:mainfrom
Topologized:fix_mma_pipeline_2
Open

fix MMA Pipeline theoretical race#190
Topologized wants to merge 1 commit intoHazyResearch:mainfrom
Topologized:fix_mma_pipeline_2

Conversation

@Topologized
Copy link
Copy Markdown

@Topologized Topologized commented Mar 29, 2026

The current B200 MMA pipeline synchronization mechanism has a theoretical bug with MMA_PIPE_DEPTH > 1 and CLC_PIPE_DEPTH > 1, where the MMA warp can arrive twice before the epilogue warp(s) wait for the inputs to arrive. Although the work of a MMA is typically quite large compared to the epilogue so that this race is unlikely to occur in practice, one can reliably reproduce the bug by a configuration such as:

    run_benchmark<config<256, 256,  64, 8, true, 6, 8>>(32768, 8192, /*K=*/64, ncu);

If one then sets CLC_PIPE_DEPTH=2 and runs compute-sanitizer --tool synccheck, errors are reported. A reliable crash will happen with CLC_PIPE_DEPTH=4 without compute-sanitizer. If CLC_PIPE_DEPTH=1, the schedule_arrived mbarrier acts as a synchronization mechanism to not cause this bug.

Although mbarriers are cheap in terms of shared memory usage, a wontfix could be acceptable if appropriate comments are added to warn the reader of this potential bug.

I tested the bf16 and fp8 normal gemms on my B200 machine; I haven't tested the ag_gemm patches, but it should work (I can look into setting the parallel GEMM up on my machine if required).

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.

1 participant