llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-llvm-ir @llvm/pr-subscribers-backend-amdgpu Author: Sameer Sahasrabuddhe (ssahasra) <details> <summary>Changes</summary> Asynchronous operations are memory transfers (usually between the global memory and LDS) that are completed independently at an unspecified scope. A thread that requests one or more asynchronous transfers can use async marks to track their completion. The thread waits for each mark to be completed, which indicates that requests initiated in program order before this mark have also completed. For now, we implement asyncmark/wait operations on pre-GFX12 architectures that support "LDS DMA" operations. Future work will extend support to GFX12Plus architectures that support "true" async operations. This is part of a stack split out from #<!-- -->173259 Co-authored-by: Ryan Mitchell ryan.mitchell@<!-- -->amd.com Fixes: SWDEV-521121 --- Patch is 70.17 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/180467.diff 13 Files Affected: - (modified) clang/include/clang/Basic/BuiltinsAMDGPU.td (+8) - (added) clang/test/CodeGenOpenCL/builtins-amdgcn-asyncmark-errs.cl (+7) - (added) clang/test/CodeGenOpenCL/builtins-amdgcn-asyncmark.cl (+16) - (modified) llvm/include/llvm/IR/IntrinsicsAMDGPU.td (+9) - (modified) llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp (+6) - (modified) llvm/lib/Target/AMDGPU/AMDGPUMCInstLower.cpp (+15-1) - (modified) llvm/lib/Target/AMDGPU/SIInsertWaitcnts.cpp (+268-12) - (modified) llvm/lib/Target/AMDGPU/SIInstrInfo.h (+4-4) - (modified) llvm/lib/Target/AMDGPU/SOPInstructions.td (+11-1) - (modified) llvm/test/CodeGen/AMDGPU/async-buffer-loads.ll (+34-12) - (added) llvm/test/CodeGen/AMDGPU/asyncmark-err.ll (+19) - (added) llvm/test/CodeGen/AMDGPU/asyncmark-max-pregfx12.ll (+279) - (modified) llvm/test/CodeGen/AMDGPU/asyncmark-pregfx12.ll (+188-69) ``````````diff diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td index 7eb6cece7c55b..507db501267a0 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.td +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td @@ -194,6 +194,14 @@ def __builtin_amdgcn_raw_ptr_buffer_load_async_lds : AMDGPUBuiltin<"void(__amdgp def __builtin_amdgcn_struct_ptr_buffer_load_lds : AMDGPUBuiltin<"void(__amdgpu_buffer_rsrc_t, void address_space<3> *, _Constant unsigned int, int, int, int, _Constant int, _Constant int)", [], "vmem-to-lds-load-insts">; def __builtin_amdgcn_struct_ptr_buffer_load_async_lds : AMDGPUBuiltin<"void(__amdgpu_buffer_rsrc_t, void address_space<3> *, _Constant unsigned int, int, int, int, _Constant int, _Constant int)", [], "vmem-to-lds-load-insts">; +//===----------------------------------------------------------------------===// +// Async mark builtins. +//===----------------------------------------------------------------------===// + +// FIXME: Not supported on GFX12 yet. Will need a new feature when we do. +def __builtin_amdgcn_asyncmark : AMDGPUBuiltin<"void()", [], "vmem-to-lds-load-insts">; +def __builtin_amdgcn_wait_asyncmark : AMDGPUBuiltin<"void(_Constant unsigned short)", [], "vmem-to-lds-load-insts">; + //===----------------------------------------------------------------------===// // Ballot builtins. //===----------------------------------------------------------------------===// diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-asyncmark-errs.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-asyncmark-errs.cl new file mode 100644 index 0000000000000..7d4a141fbde6e --- /dev/null +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-asyncmark-errs.cl @@ -0,0 +1,7 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -O0 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -target-cpu gfx1250 -verify -S -o - %s + +void test_feature() { + __builtin_amdgcn_asyncmark(); // expected-error{{'__builtin_amdgcn_asyncmark' needs target feature vmem-to-lds-load-insts}} + __builtin_amdgcn_wait_asyncmark(0); // expected-error{{'__builtin_amdgcn_wait_asyncmark' needs target feature vmem-to-lds-load-insts}} +} diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-asyncmark.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-asyncmark.cl new file mode 100644 index 0000000000000..976ae3cea5d6d --- /dev/null +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-asyncmark.cl @@ -0,0 +1,16 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py +// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu gfx900 -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu gfx942 -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -emit-llvm -o - %s | FileCheck %s +// REQUIRES: amdgpu-registered-target + +// CHECK-LABEL: @test_invocation( +// CHECK-NEXT: entry: +// CHECK-NEXT: call void @llvm.amdgcn.asyncmark() +// CHECK-NEXT: call void @llvm.amdgcn.wait.asyncmark(i16 0) +// CHECK-NEXT: ret void +// +void test_invocation() { + __builtin_amdgcn_asyncmark(); + __builtin_amdgcn_wait_asyncmark(0); +} diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td index ec887ab0bdfad..3d3631b5d92e8 100644 --- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td +++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td @@ -2826,6 +2826,15 @@ def int_amdgcn_global_load_async_lds : AMDGPUGlobalLoadLDS, ClangBuiltin<"__buil def int_amdgcn_pops_exiting_wave_id : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrHasSideEffects]>; +// Sets a marker in the stream of async requests. Modelled as InaccessibleMem. +def int_amdgcn_asyncmark : ClangBuiltin<"__builtin_amdgcn_asyncmark">, + Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects]>; + +// Waits until the Nth previous marker is completed, if it exists. +def int_amdgcn_wait_asyncmark : + ClangBuiltin<"__builtin_amdgcn_wait_asyncmark">, + Intrinsic<[], [llvm_i16_ty], [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects]>; + //===----------------------------------------------------------------------===// // GFX10 Intrinsics //===----------------------------------------------------------------------===// diff --git a/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp b/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp index f2e5be546a11d..e1fe181bd0140 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp @@ -2393,6 +2393,12 @@ bool AMDGPUInstructionSelector::selectG_INTRINSIC_W_SIDE_EFFECTS( case Intrinsic::amdgcn_global_load_lds: case Intrinsic::amdgcn_global_load_async_lds: return selectGlobalLoadLds(I); + case Intrinsic::amdgcn_asyncmark: + case Intrinsic::amdgcn_wait_asyncmark: + // FIXME: Not supported on GFX12 yet. Will need a new feature when we do. + if (!Subtarget->hasVMemToLDSLoad()) + return false; + break; case Intrinsic::amdgcn_exp_compr: if (!STI.hasCompressedExport()) { Function &F = I.getMF()->getFunction(); diff --git a/llvm/lib/Target/AMDGPU/AMDGPUMCInstLower.cpp b/llvm/lib/Target/AMDGPU/AMDGPUMCInstLower.cpp index 99c1ab8d379d5..fc408aa30dd87 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUMCInstLower.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUMCInstLower.cpp @@ -347,7 +347,7 @@ void AMDGPUAsmPrinter::emitInstruction(const MachineInstr *MI) { } } else { // We don't want these pseudo instructions encoded. They are - // placeholder terminator instructions and should only be printed as + // placeholder instructions and should only be printed as // comments. if (MI->getOpcode() == AMDGPU::SI_RETURN_TO_EPILOG) { if (isVerbose()) @@ -361,6 +361,20 @@ void AMDGPUAsmPrinter::emitInstruction(const MachineInstr *MI) { return; } + if (MI->getOpcode() == AMDGPU::ASYNCMARK) { + if (isVerbose()) + OutStreamer->emitRawComment(" asyncmark"); + return; + } + + if (MI->getOpcode() == AMDGPU::WAIT_ASYNCMARK) { + if (isVerbose()) { + OutStreamer->emitRawComment(" wait_asyncmark(" + + Twine(MI->getOperand(0).getImm()) + ")"); + } + return; + } + if (MI->getOpcode() == AMDGPU::SCHED_BARRIER) { if (isVerbose()) { std::string HexString; diff --git a/llvm/lib/Target/AMDGPU/SIInsertWaitcnts.cpp b/llvm/lib/Target/AMDGPU/SIInsertWaitcnts.cpp index 2014a415f3121..84dec5e7ccfac 100644 --- a/llvm/lib/Target/AMDGPU/SIInsertWaitcnts.cpp +++ b/llvm/lib/Target/AMDGPU/SIInsertWaitcnts.cpp @@ -710,6 +710,24 @@ class SIInsertWaitcnts { std::optional<WaitEventType> getExpertSchedulingEventType(const MachineInstr &Inst) const; + bool isAsync(const MachineInstr &MI) const { + if (!SIInstrInfo::isLDSDMA(MI)) + return false; + if (SIInstrInfo::usesASYNC_CNT(MI)) + return true; + const MachineOperand *Async = + TII->getNamedOperand(MI, AMDGPU::OpName::IsAsync); + return Async && (Async->getImm()); + } + + bool isNonAsyncLdsDmaWrite(const MachineInstr &MI) const { + return SIInstrInfo::mayWriteLDSThroughDMA(MI) && !isAsync(MI); + } + + bool isAsyncLdsDmaWrite(const MachineInstr &MI) const { + return SIInstrInfo::mayWriteLDSThroughDMA(MI) && isAsync(MI); + } + bool isVmemAccess(const MachineInstr &MI) const; bool generateWaitcntInstBefore(MachineInstr &MI, WaitcntBrackets &ScoreBrackets, @@ -838,11 +856,13 @@ class WaitcntBrackets { AMDGPU::Waitcnt &Wait) const; void determineWaitForLDSDMA(InstCounterType T, VMEMID TID, AMDGPU::Waitcnt &Wait) const; + AMDGPU::Waitcnt determineAsyncWait(unsigned N); void tryClearSCCWriteEvent(MachineInstr *Inst); void applyWaitcnt(const AMDGPU::Waitcnt &Wait); void applyWaitcnt(InstCounterType T, unsigned Count); void updateByEvent(WaitEventType E, MachineInstr &MI); + void recordAsyncMark(MachineInstr &MI); bool hasPendingEvent() const { return !PendingEvents.empty(); } bool hasPendingEvent(WaitEventType E) const { @@ -934,11 +954,15 @@ class WaitcntBrackets { unsigned OtherShift; }; + using CounterValueArray = std::array<unsigned, NUM_INST_CNTS>; + void determineWaitForScore(InstCounterType T, unsigned Score, AMDGPU::Waitcnt &Wait) const; static bool mergeScore(const MergeInfo &M, unsigned &Score, unsigned OtherScore); + bool mergeAsyncMarks(ArrayRef<MergeInfo> MergeInfos, + ArrayRef<CounterValueArray> OtherMarks); iterator_range<MCRegUnitIterator> regunits(MCPhysReg Reg) const { assert(Reg != AMDGPU::SCC && "Shouldn't be used on SCC"); @@ -1015,8 +1039,8 @@ class WaitcntBrackets { // TODO: Could we track SCC alongside SGPRs so it's not longer a special case? struct VMEMInfo { - // Scores for all instruction counters. - std::array<unsigned, NUM_INST_CNTS> Scores = {0}; + // Scores for all instruction counters. Zero-initialized. + CounterValueArray Scores{}; // Bitmask of the VmemTypes of VMEM instructions for this VGPR. unsigned VMEMTypes = 0; @@ -1044,6 +1068,14 @@ class WaitcntBrackets { // Store representative LDS DMA operations. The only useful info here is // alias info. One store is kept per unique AAInfo. SmallVector<const MachineInstr *> LDSDMAStores; + + // State of all counters at each async mark encountered so far. + SmallVector<CounterValueArray> AsyncMarks; + static constexpr unsigned MaxAsyncMarks = 16; + + // Track the upper bound score for async operations that are not part of a + // mark yet. Initialized to all zeros. + CounterValueArray AsyncScore{}; }; class SIInsertWaitcntsLegacy : public MachineFunctionPass { @@ -1245,7 +1277,7 @@ void WaitcntBrackets::updateByEvent(WaitEventType E, MachineInstr &Inst) { setScoreByOperand(Op, T, CurrScore); } if (Inst.mayStore() && - (TII->isDS(Inst) || TII->mayWriteLDSThroughDMA(Inst))) { + (TII->isDS(Inst) || Context->isNonAsyncLdsDmaWrite(Inst))) { // MUBUF and FLAT LDS DMA operations need a wait on vmcnt before LDS // written can be accessed. A load from LDS to VMEM does not need a wait. // @@ -1289,6 +1321,14 @@ void WaitcntBrackets::updateByEvent(WaitEventType E, MachineInstr &Inst) { setVMemScore(LDSDMA_BEGIN + Slot, T, CurrScore); } + // FIXME: Not supported on GFX12 yet. Newer async operations use other + // counters too, so will need a map from instruction or event types to + // counter types. + if (Context->isAsyncLdsDmaWrite(Inst) && T == LOAD_CNT) { + assert(!SIInstrInfo::usesASYNC_CNT(Inst)); + AsyncScore[T] = CurrScore; + } + if (SIInstrInfo::isSBarrierSCCWrite(Inst.getOpcode())) { setRegScore(AMDGPU::SCC, T, CurrScore); PendingSCCWrite = &Inst; @@ -1296,13 +1336,28 @@ void WaitcntBrackets::updateByEvent(WaitEventType E, MachineInstr &Inst) { } } +void WaitcntBrackets::recordAsyncMark(MachineInstr &Inst) { + // In the absence of loops, AsyncMarks can grow linearly with the program + // until we encounter an ASYNCMARK_WAIT. We could drop the oldest mark above a + // limit every time we push a new mark, but that seems like unnecessary work + // in practical cases. We do separately truncate the array when processing a + // loop, which should be sufficient. + AsyncMarks.push_back(AsyncScore); + AsyncScore = {}; + LLVM_DEBUG({ + dbgs() << "recordAsyncMark:\n" << Inst; + for (const auto &Mark : AsyncMarks) { + llvm::interleaveComma(Mark, dbgs()); + dbgs() << '\n'; + } + }); +} + void WaitcntBrackets::print(raw_ostream &OS) const { const GCNSubtarget *ST = Context->ST; - OS << '\n'; for (auto T : inst_counter_types(Context->MaxCounter)) { unsigned SR = getScoreRange(T); - switch (T) { case LOAD_CNT: OS << " " << (ST->hasExtendedWaitCounts() ? "LOAD" : "VM") << "_CNT(" @@ -1395,6 +1450,53 @@ void WaitcntBrackets::print(raw_ostream &OS) const { } OS << '\n'; + OS << "Async score: "; + if (AsyncScore.empty()) + OS << "none"; + else + llvm::interleaveComma(AsyncScore, OS); + OS << '\n'; + + OS << "Async marks: " << AsyncMarks.size() << '\n'; + + for (const auto &Mark : AsyncMarks) { + for (auto T : inst_counter_types()) { + unsigned MarkedScore = Mark[T]; + switch (T) { + case LOAD_CNT: + OS << " " << (ST->hasExtendedWaitCounts() ? "LOAD" : "VM") + << "_CNT: " << MarkedScore; + break; + case DS_CNT: + OS << " " << (ST->hasExtendedWaitCounts() ? "DS" : "LGKM") + << "_CNT: " << MarkedScore; + break; + case EXP_CNT: + OS << " EXP_CNT: " << MarkedScore; + break; + case STORE_CNT: + OS << " " << (ST->hasExtendedWaitCounts() ? "STORE" : "VS") + << "_CNT: " << MarkedScore; + break; + case SAMPLE_CNT: + OS << " SAMPLE_CNT: " << MarkedScore; + break; + case BVH_CNT: + OS << " BVH_CNT: " << MarkedScore; + break; + case KM_CNT: + OS << " KM_CNT: " << MarkedScore; + break; + case X_CNT: + OS << " X_CNT: " << MarkedScore; + break; + default: + OS << " UNKNOWN: " << MarkedScore; + break; + } + } + OS << '\n'; + } OS << '\n'; } @@ -1496,6 +1598,49 @@ void WaitcntBrackets::determineWaitForScore(InstCounterType T, } } +AMDGPU::Waitcnt WaitcntBrackets::determineAsyncWait(unsigned N) { + LLVM_DEBUG({ + dbgs() << "Need " << N << " async marks. Found " << AsyncMarks.size() + << ":\n"; + for (const auto &Mark : AsyncMarks) { + llvm::interleaveComma(Mark, dbgs()); + dbgs() << '\n'; + } + }); + + AMDGPU::Waitcnt Wait; + if (AsyncMarks.size() == MaxAsyncMarks) { + // Enforcing MaxAsyncMarks here is unnecessary work because the size of + // MaxAsyncMarks is linear when traversing straightline code. But we do + // need to check if truncation may have occured at a merge, and adjust N + // to ensure that a wait is generated. + LLVM_DEBUG(dbgs() << "Possible truncation. Ensuring a non-trivial wait.\n"); + N = std::min(N, (unsigned)MaxAsyncMarks - 1); + } + + if (AsyncMarks.size() <= N) { + LLVM_DEBUG(dbgs() << "No additional wait for async mark.\n"); + return Wait; + } + + size_t MarkIndex = AsyncMarks.size() - N - 1; + const auto &RequiredMark = AsyncMarks[MarkIndex]; + for (InstCounterType T : inst_counter_types()) + determineWaitForScore(T, RequiredMark[T], Wait); + + // Immediately remove the waited mark and all older ones + // This happens BEFORE the wait is actually inserted, which is fine + // because we've already extracted the wait requirements + LLVM_DEBUG({ + dbgs() << "Removing " << (MarkIndex + 1) + << " async marks after determining wait\n"; + }); + AsyncMarks.erase(AsyncMarks.begin(), AsyncMarks.begin() + MarkIndex + 1); + + LLVM_DEBUG(dbgs() << "Waits to add: " << Wait); + return Wait; +} + void WaitcntBrackets::determineWaitForPhysReg(InstCounterType T, MCPhysReg Reg, AMDGPU::Waitcnt &Wait) const { if (Reg == AMDGPU::SCC) { @@ -1723,6 +1868,11 @@ bool WaitcntGeneratorPreGFX12::applyPreexistingWaitcnt( // possibility in an articial MIR test since such a situation cannot be // recreated by running the memory legalizer. II.eraseFromParent(); + } else if (Opcode == AMDGPU::WAIT_ASYNCMARK) { + unsigned N = II.getOperand(0).getImm(); + LLVM_DEBUG(dbgs() << "Processing WAIT_ASYNCMARK: " << II << '\n';); + AMDGPU::Waitcnt OldWait = ScoreBrackets.determineAsyncWait(N); + Wait = Wait.combined(OldWait); } else { assert(Opcode == AMDGPU::S_WAITCNT_VSCNT); assert(II.getOperand(0).getReg() == AMDGPU::SGPR_NULL); @@ -1976,6 +2126,8 @@ bool WaitcntGeneratorGFX12Plus::applyPreexistingWaitcnt( // LDS, so no work required here yet. II.eraseFromParent(); continue; + } else if (Opcode == AMDGPU::WAIT_ASYNCMARK) { + reportFatalUsageError("WAIT_ASYNCMARK is not ready for GFX12 yet"); } else { std::optional<InstCounterType> CT = counterTypeForInstr(Opcode); assert(CT.has_value()); @@ -2300,6 +2452,7 @@ bool WaitcntGeneratorGFX12Plus::createNewWaitcnt( bool SIInsertWaitcnts::generateWaitcntInstBefore( MachineInstr &MI, WaitcntBrackets &ScoreBrackets, MachineInstr *OldWaitcntInstr, PreheaderFlushFlags FlushFlags) { + LLVM_DEBUG(dbgs() << "\n*** GenerateWaitcntInstBefore: "; MI.print(dbgs());); setForceEmitWaitcnt(); assert(!MI.isMetaInstruction()); @@ -2854,6 +3007,84 @@ bool WaitcntBrackets::mergeScore(const MergeInfo &M, unsigned &Score, return OtherShifted > MyShifted; } +bool WaitcntBrackets::mergeAsyncMarks(ArrayRef<MergeInfo> MergeInfos, + ArrayRef<CounterValueArray> OtherMarks) { + bool StrictDom = false; + + LLVM_DEBUG(dbgs() << "Merging async marks ..."); + // Early exit: both empty + if (AsyncMarks.empty() && OtherMarks.empty()) { + LLVM_DEBUG(dbgs() << " nothing to merge\n"); + return false; + } + LLVM_DEBUG(dbgs() << '\n'); + + // Determine maximum length needed after merging + auto MaxSize = (unsigned)std::max(AsyncMarks.size(), OtherMarks.size()); + + // For each backedge in isolation, the algorithm reachs a fixed point after + // the first call to merge(). This is unchanged even with the AsyncMarks + // array because we call mergeScore just like the other cases. + // + // But in the rare pathological case, a nest of loops that pushes marks + // without waiting on any mark can cause AsyncMarks to grow very large. We cap + // it to a reasonable limit. We can tune this later or potentially introduce a + // user option to control the value. + MaxSize = std::min(MaxSize, MaxAsyncMarks); + + // Keep only the most recent marks within our limit. + if (AsyncMarks.size() > MaxSize) + AsyncMarks.erase(AsyncMarks.begin(), + AsyncMarks.begin() + (AsyncMarks.size() - MaxSize)); + + // Pad with zero-filled marks if our list is shorter. Zero represents "no + // pending async operations at this checkpoint" and acts as the identity + // element for max() during merging. We pad at the beginning since the marks + // need to be aligned in most-recent order. + CounterValueArray ZeroMark{}; + AsyncMarks.insert(AsyncMarks.begin(), MaxSize - AsyncMarks.size(), ZeroMark); + + LLVM_DEBUG({ + dbgs() << "Before merge:\n"; + for (const auto &Mark : AsyncMarks) { + llvm::interleaveComma(Mark, dbgs()); + dbgs() << '\n'; + } + }); + + LLVM_DEBUG({ + dbgs() << "Other marks:\n"; + for (const auto &Mark : OtherMarks) { + llvm::interleaveComma(Mark, dbgs()); + dbgs() << '\n'; + } + }); + + // Merge element-wise using the existing mergeScore function and the + // appropriate MergeInfo for each counter type. Iterate only while we have + // elements in both vectors. + unsigned OtherSize = OtherMarks.size(); + unsigned OurSize = AsyncMarks.size(); + unsigned MergeCount = std::min(OtherSize, OurSize); + assert(OurSize == MaxSize); + for (unsigned Idx = 1; Idx <= MergeCount; ++Idx) { + for (auto T : inst_counter_types(Context->MaxCounter)) { + StrictDom |= mergeScore(MergeInfos[T], AsyncMarks[OurSize - Idx][T], + OtherMarks[OtherSize - Idx][T]); + } + } + + LLVM_DEBUG({ + dbgs() << "After merge:\n"; + for (const auto &Mark : AsyncMarks) { + llvm::interleaveComma(Mark, dbgs()); + dbgs() << '\n'; + } + }); + + return StrictDom; +} + /// Merge the pending events and associater score brackets of \p Other into /// this brackets status. /// @@ -2869,6 +3100,9 @@ bool WaitcntBrackets::merge(const WaitcntBrackets &Other) { for (auto K : Other.SGPRs.keys()) SGPRs.try_emplace(K); + // Array to store MergeInfo for each counter type + MergeInfo MergeInfos[NUM_INST_CNTS]; + for (auto T : inst_counter_types(Context->MaxCounter)) { // Merge event flags for this counter const WaitEventSet &EventsForT = Context->get... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/180467 _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
