Skip to content

Commit 2263431

Browse files
authored
[WS] Add a missing fence after local_store for MMAv5 consumers (#8317)
1 parent 11c876b commit 2263431

File tree

4 files changed

+164
-3
lines changed

4 files changed

+164
-3
lines changed

lib/Dialect/TritonInstrument/Transforms/ConcurrencySanitizer.cpp

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -568,7 +568,8 @@ class ConcurrencySanitizerPass
568568
}
569569
if (auto wgmmaOp = dyn_cast<ttng::WarpGroupDotOp>(op)) {
570570
if (wgmmaOp.getIsAsync() == true) {
571-
info = {.trackingKind = MemEffectsOpInfo::TrackingKind::wgmmaCommit};
571+
info = {.trackingKind = MemEffectsOpInfo::TrackingKind::wgmmaCommit,
572+
.barriers = {}};
572573
if (isa<ttg::SharedEncodingTrait>(
573574
wgmmaOp.getA().getType().getEncoding())) {
574575
info->operandEffects.emplace_back(MemEffectsOpInfo::Effects{

0 commit comments

Comments
 (0)