Skip to content

Commit 5745035

Browse files
anmyachevanmyachev
authored
Do not use C++20 designed initializers in Transforms/ConcurrencySanitizer.cpp (#8455)
Signed-off-by: Anatoly Myachev <[email protected]>
1 parent 35bb255 commit 5745035

File tree

1 file changed

+40
-36
lines changed

1 file changed

+40
-36
lines changed

‎lib/Dialect/TritonInstrument/Transforms/ConcurrencySanitizer.cpp‎

Lines changed: 40 additions & 36 deletions
Original file line numberDiff line numberDiff line change
@@ -459,79 +459,83 @@ class ConcurrencySanitizerPass
459459
info->pred = copyOp.getPred();
460460
info->barriers.push_back({copyOp.getBarrier(), nullptr, 1});
461461
info->operandEffects.push_back(
462-
{.rw = MemEffectsOpInfo::Effects::RW::Write,
463-
.buf = copyOp.getResult()});
462+
{/*.rw =*/MemEffectsOpInfo::Effects::RW::Write,
463+
/*.buf =*/copyOp.getResult()});
464464
}
465465
if (auto storeOp = dyn_cast<ttng::AsyncTMACopyLocalToGlobalOp>(op)) {
466466
info.emplace();
467467
info->trackingKind = MemEffectsOpInfo::TrackingKind::None;
468468
info->operandEffects.push_back(
469-
{.rw = MemEffectsOpInfo::Effects::RW::Read, .buf = storeOp.getSrc()});
469+
{/*.rw =*/MemEffectsOpInfo::Effects::RW::Read,
470+
/*.buf =*/storeOp.getSrc()});
470471
}
471472
if (auto gatherOp = dyn_cast<ttng::AsyncTMAGatherOp>(op)) {
472473
info.emplace();
473474
info->trackingKind = MemEffectsOpInfo::TrackingKind::Barrier;
474475
info->pred = gatherOp.getPred();
475476
info->barriers.push_back({gatherOp.getBarrier(), nullptr, 1});
476477
info->operandEffects.push_back(
477-
{.rw = MemEffectsOpInfo::Effects::RW::Write,
478-
.buf = gatherOp.getResult()});
478+
{/*.rw =*/MemEffectsOpInfo::Effects::RW::Write,
479+
/*.buf =*/gatherOp.getResult()});
479480
}
480481
if (auto scatterOp = dyn_cast<ttng::AsyncTMAScatterOp>(op)) {
481482
info.emplace();
482483
info->trackingKind = MemEffectsOpInfo::TrackingKind::None;
483-
info->operandEffects.push_back({.rw = MemEffectsOpInfo::Effects::RW::Read,
484-
.buf = scatterOp.getSrc()});
484+
info->operandEffects.push_back(
485+
{/*.rw =*/MemEffectsOpInfo::Effects::RW::Read,
486+
/*.buf =*/scatterOp.getSrc()});
485487
}
486488
if (auto copyOp = dyn_cast<ttg::AsyncCopyGlobalToLocalOp>(op)) {
487489
info.emplace();
488490
info->trackingKind = MemEffectsOpInfo::TrackingKind::asyncCpCommit;
489491
info->operandEffects.push_back(
490-
{.rw = MemEffectsOpInfo::Effects::RW::Write,
491-
.buf = copyOp.getResult()});
492+
{/*.rw =*/MemEffectsOpInfo::Effects::RW::Write,
493+
/*.buf =*/copyOp.getResult()});
492494
}
493495
if (auto loadOp = dyn_cast<ttg::LocalLoadOp>(op)) {
494496
info.emplace();
495497
info->trackingKind = MemEffectsOpInfo::TrackingKind::Barrier;
496498
info->operandEffects.push_back(
497-
{.rw = MemEffectsOpInfo::Effects::RW::Read, .buf = loadOp.getSrc()});
499+
{/*.rw =*/MemEffectsOpInfo::Effects::RW::Read,
500+
/*.buf =*/loadOp.getSrc()});
498501
}
499502
if (auto storeOp = dyn_cast<ttg::LocalStoreOp>(op)) {
500503
info.emplace();
501504
info->trackingKind = MemEffectsOpInfo::TrackingKind::Barrier;
502505
info->operandEffects.push_back(
503-
{.rw = MemEffectsOpInfo::Effects::RW::Write,
504-
.buf = storeOp.getDst()});
506+
{/*.rw =*/MemEffectsOpInfo::Effects::RW::Write,
507+
/*.buf =*/storeOp.getDst()});
505508
}
506509
if (auto allocOp = dyn_cast<ttg::LocalAllocOp>(op)) {
507510
if (allocOp.getSrc()) {
508511
info.emplace();
509512
info->trackingKind = MemEffectsOpInfo::TrackingKind::Barrier;
510513
info->operandEffects.push_back(
511-
{.rw = MemEffectsOpInfo::Effects::RW::Write,
512-
.buf = allocOp.getResult()});
514+
{/*.rw =*/MemEffectsOpInfo::Effects::RW::Write,
515+
/*.buf =*/allocOp.getResult()});
513516
}
514517
}
515518
if (auto loadOp = dyn_cast<ttng::TMEMLoadOp>(op)) {
516519
info.emplace();
517520
info->trackingKind = MemEffectsOpInfo::TrackingKind::Barrier;
518521
info->operandEffects.push_back(
519-
{.rw = MemEffectsOpInfo::Effects::RW::Read, .buf = loadOp.getSrc()});
522+
{/*.rw =*/MemEffectsOpInfo::Effects::RW::Read,
523+
/*.buf =*/loadOp.getSrc()});
520524
}
521525
if (auto storeOp = dyn_cast<ttng::TMEMStoreOp>(op)) {
522526
info.emplace();
523527
info->trackingKind = MemEffectsOpInfo::TrackingKind::Barrier;
524528
info->operandEffects.push_back(
525-
{.rw = MemEffectsOpInfo::Effects::RW::Write,
526-
.buf = storeOp.getDst()});
529+
{/*.rw =*/MemEffectsOpInfo::Effects::RW::Write,
530+
/*.buf =*/storeOp.getDst()});
527531
}
528532
if (auto allocOp = dyn_cast<ttng::TMEMAllocOp>(op)) {
529533
if (allocOp.getSrc()) {
530534
info.emplace();
531535
info->trackingKind = MemEffectsOpInfo::TrackingKind::Barrier;
532536
info->operandEffects.push_back(
533-
{.rw = MemEffectsOpInfo::Effects::RW::Write,
534-
.buf = allocOp.getResult()});
537+
{/*.rw =*/MemEffectsOpInfo::Effects::RW::Write,
538+
/*.buf =*/allocOp.getResult()});
535539
}
536540
}
537541
if (auto mmav5Op = dyn_cast<ttng::TCGen5MMAOp>(op)) {
@@ -542,16 +546,15 @@ class ConcurrencySanitizerPass
542546
llvm::zip(mmav5Op.getBarriers(), mmav5Op.getBarrierPreds())) {
543547
info->barriers.push_back({barrier, barrierPred, 1});
544548
}
545-
info->operandEffects.push_back({.rw = MemEffectsOpInfo::Effects::RW::Read,
546-
.buf = mmav5Op.getA(),
547-
.operandName = "A"});
548-
info->operandEffects.push_back({.rw = MemEffectsOpInfo::Effects::RW::Read,
549-
.buf = mmav5Op.getB(),
550-
.operandName = "B"});
551549
info->operandEffects.push_back(
552-
{.rw = MemEffectsOpInfo::Effects::RW::Write,
553-
.buf = mmav5Op.getAccumulator(),
554-
.operandName = "Acc"});
550+
{/*.rw =*/MemEffectsOpInfo::Effects::RW::Read,
551+
/*.buf =*/mmav5Op.getA(), /*.operandName =*/"A"});
552+
info->operandEffects.push_back(
553+
{/*.rw =*/MemEffectsOpInfo::Effects::RW::Read,
554+
/*.buf =*/mmav5Op.getB(), /*.operandName =*/"B"});
555+
info->operandEffects.push_back(
556+
{/*.rw =*/MemEffectsOpInfo::Effects::RW::Write,
557+
/*.buf =*/mmav5Op.getAccumulator(), /*.operandName =*/"Acc"});
555558
}
556559
if (auto commitOp = dyn_cast<ttng::TCGen5CommitOp>(op)) {
557560
info.emplace();
@@ -568,21 +571,22 @@ class ConcurrencySanitizerPass
568571
}
569572
if (auto wgmmaOp = dyn_cast<ttng::WarpGroupDotOp>(op)) {
570573
if (wgmmaOp.getIsAsync() == true) {
571-
info = {.trackingKind = MemEffectsOpInfo::TrackingKind::wgmmaCommit,
572-
.barriers = {}};
574+
info.emplace();
575+
info->trackingKind = MemEffectsOpInfo::TrackingKind::wgmmaCommit;
576+
info->barriers = {};
573577
if (isa<ttg::SharedEncodingTrait>(
574578
wgmmaOp.getA().getType().getEncoding())) {
575579
info->operandEffects.emplace_back(MemEffectsOpInfo::Effects{
576-
.rw = MemEffectsOpInfo::Effects::RW::Read,
577-
.buf = wgmmaOp.getA(),
578-
.operandName = "A"});
580+
/*.rw =*/MemEffectsOpInfo::Effects::RW::Read,
581+
/*.buf =*/wgmmaOp.getA(),
582+
/*.operandName =*/"A"});
579583
}
580584
if (isa<ttg::SharedEncodingTrait>(
581585
wgmmaOp.getB().getType().getEncoding())) {
582586
info->operandEffects.emplace_back(MemEffectsOpInfo::Effects{
583-
.rw = MemEffectsOpInfo::Effects::RW::Read,
584-
.buf = wgmmaOp.getB(),
585-
.operandName = "B"});
587+
/*.rw =*/MemEffectsOpInfo::Effects::RW::Read,
588+
/*.buf =*/wgmmaOp.getB(),
589+
/*.operandName =*/"B"});
586590
}
587591
}
588592
}

0 commit comments

Comments
 (0)