From 96cd3796df7f493dec2c93222c5d3ec5d38396d3 Mon Sep 17 00:00:00 2001 From: Steven Inouye Date: Sun, 12 Apr 2026 13:37:40 -0700 Subject: [PATCH 01/30] Added: single-thread implementation --- include/opt-sched/Scheduler/aco.h | 5 ++++ lib/Scheduler/aco.hip.cpp | 43 +++++++++++++++++++++++++++++-- 2 files changed, 46 insertions(+), 2 deletions(-) diff --git a/include/opt-sched/Scheduler/aco.h b/include/opt-sched/Scheduler/aco.h index 3fd688dd..b3360a0c 100644 --- a/include/opt-sched/Scheduler/aco.h +++ b/include/opt-sched/Scheduler/aco.h @@ -77,6 +77,11 @@ class ACOScheduler : public ConstrainedScheduler { void AllocDevArraysForParallelACO(); // Finds a schedule, if passed a device side schedule, use that instead // of creating a new one + + + InstSchedule *FindManyCPUSchedule(InstCount RPTarget); + + __host__ __device__ InstSchedule *FindOneSchedule(InstCount RPTarget, InstSchedule *dev_schedule = NULL, int kernelNum = -1); diff --git a/lib/Scheduler/aco.hip.cpp b/lib/Scheduler/aco.hip.cpp index 09aa9b83..752f2e45 100644 --- a/lib/Scheduler/aco.hip.cpp +++ b/lib/Scheduler/aco.hip.cpp @@ -16,6 +16,9 @@ #include #include +#include +#include + using namespace llvm::opt_sched; namespace cg = cooperative_groups; @@ -51,6 +54,9 @@ double RandDouble(double min, double max) { //#define DECAY_FACTOR 0.5 //#endif +#define RUN_PCPU 1 +#define NO_CPU_THREADS 1 + ACOScheduler::ACOScheduler(DataDepGraph *dataDepGraph, MachineModel *machineModel, InstCount upperBound, SchedPriorities priorities1, SchedPriorities priorities2, bool vrfySched, @@ -1788,12 +1794,19 @@ FUNC_RESULT ACOScheduler::FindSchedule(InstSchedule *schedule_out, std::unordered_map schedMap; int diffSchedCount = 0; #endif - + while (noImprovement < noImprovementMax) { iterations++; iterationBest = nullptr; for (int i = 0; i < numThreads_; i++) { - InstSchedule *schedule = FindOneSchedule(RPTarget, NULL); + InstSchedule *schedule; + if (!RUN_PCPU) { + schedule = FindOneSchedule(RPTarget, NULL); + } else { + schedule = FindManyCPUSchedule(RPTarget); + i+=(NO_CPU_THREADS-1); + } + #ifdef CHECK_DIFFERENT_SCHEDULES // check if schedule is in Map @@ -1871,6 +1884,7 @@ FUNC_RESULT ACOScheduler::FindSchedule(InstSchedule *schedule_out, #if USE_ACS UpdatePheromone(bestSchedule, false); #endif + } Logger::Info("%d ants terminated early", numAntsTerminated_); #ifdef CHECK_DIFFERENT_SCHEDULES @@ -1889,6 +1903,7 @@ FUNC_RESULT ACOScheduler::FindSchedule(InstSchedule *schedule_out, if (bestSchedule != InitialSchedule) delete bestSchedule; printf("Occupancy Target %d had %d iterations\n",j ,iterations); + } // End run on CPU } if (!use_dev_ACO || count_ < REGION_MIN_SIZE) @@ -2408,3 +2423,27 @@ void ACOScheduler::FreeDevicePointers(bool IsSecondPass) { hipFree(dev_kHelper2); hipFree(pheromone_.elmnts_); } + +InstSchedule *ACOScheduler::FindManyCPUSchedule(InstCount RPTarget) { + std::vector PCPUThreads; + InstSchedule **cpuScheds = new InstSchedule*[NO_CPU_THREADS](); + + for (int i = 0; i < NO_CPU_THREADS; i++) { + PCPUThreads.emplace_back([this, cpuScheds, i, RPTarget]() { + cpuScheds[i] = FindOneSchedule(RPTarget, NULL); + }); + } + + for (auto &t : PCPUThreads) { + t.join(); + } + + //logic to find best schedule + InstSchedule *result = cpuScheds[0]; + + for (int i = 1; i < NO_CPU_THREADS; i++) { + delete cpuScheds[i]; + } + delete[] cpuScheds; + return result; +} From a75ac18306c3ee4fb2e15cb7244a4f81c992d8e6 Mon Sep 17 00:00:00 2001 From: Steven Inouye Date: Sun, 12 Apr 2026 13:51:03 -0700 Subject: [PATCH 02/30] Unconfirmed: adjustments to BBwithSpill and REgister objects. --- include/opt-sched/Scheduler/aco.h | 22 ++ include/opt-sched/Scheduler/bb_spill.h | 33 +++ include/opt-sched/Scheduler/register.h | 14 + lib/Scheduler/bb_spill.hip.cpp | 368 +++++++++++++++++++++++++ lib/Scheduler/register.hip.cpp | 24 ++ 5 files changed, 461 insertions(+) diff --git a/include/opt-sched/Scheduler/aco.h b/include/opt-sched/Scheduler/aco.h index b3360a0c..d0199de9 100644 --- a/include/opt-sched/Scheduler/aco.h +++ b/include/opt-sched/Scheduler/aco.h @@ -50,6 +50,28 @@ struct BlockDecision { int blockOccupancyNum; // range 0-(difference in occupancy from AMD's schedule) }; +struct alignas(64) ParallelCPUVars { + InstCount crntCycleNum_; + InstCount crntSlotNum_; + InstCount crntSpillCost_; + InstCount crntStepNum_; + InstCount peakSpillCost_; + InstCount totSpillCost_; + InstCount slilSpillCost_; + InstCount dynamicSlilLowerBound_; + int schduldInstCnt_; + int schduldEntryInstCnt_; + int schduldExitInstCnt_; + + // These are per-thread copies of the arrays — you'll need to alloc these + WeightedBitVector *liveRegs_; // array[regTypeCnt_] + WeightedBitVector *livePhysRegs_; // array[regTypeCnt_] + InstCount *peakRegPressures_; // array[regTypeCnt_] + InstCount *spillCosts_; // array[instCnt] + int *sumOfLiveIntervalLengths_; // array[regTypeCnt_] + SmallVector regPressures_; // or just unsigned* +}; + class ACOScheduler : public ConstrainedScheduler { public: ACOScheduler(DataDepGraph *dataDepGraph, MachineModel *machineModel, diff --git a/include/opt-sched/Scheduler/bb_spill.h b/include/opt-sched/Scheduler/bb_spill.h index a2d807b1..323f91f3 100644 --- a/include/opt-sched/Scheduler/bb_spill.h +++ b/include/opt-sched/Scheduler/bb_spill.h @@ -18,6 +18,8 @@ Last Update: Apr. 2011 #include #include +#define RUN_PARALLEL_CPU 1 + namespace llvm { namespace opt_sched { @@ -29,6 +31,7 @@ class BitVector; class BBWithSpill : public SchedRegion { private: + ParallelCPUVars *pcpu_vars_; LengthCostEnumerator *enumrtr_; InstCount crntSpillCost_; @@ -267,6 +270,36 @@ class BBWithSpill : public SchedRegion { // size_t calculateMemoryNeeded() { // return regTypeCnt_ * sizeof(WeightedBitVector) * numThreads * 2; // } + + void AllocParallelCPUVars(int numThreads); + void FreeParallelCPUVars(int numThreads); + + void PCPU_CmputCrntSpillCost_(ParallelCPUVars &pcpu); + void PCPU_UpdateSpillInfoForSchdul_(SchedInstruction *inst, + bool trackCnflcts, + ParallelCPUVars &pcpu, + int thread); + void PCPU_SchdulInst(SchedInstruction *inst, InstCount cycleNum, + InstCount slotNum, bool trackCnflcts, + ParallelCPUVars &pcpu, + int thread); + + InstCount PCPU_GetCrntSpillCost(ParallelCPUVars &pcpu); + InstCount PCPU_ReturnPeakSpillCost(ParallelCPUVars &pcpu); + InstCount PCPU_getOccupancy(ParallelCPUVars &pcpu); + bool PCPU_closeToRPConstraint(ParallelCPUVars &pcpu, int blockOccupancyNum = 0); + bool PCPU_IsRPHigh(int regType, ParallelCPUVars &pcpu) const; + + InstCount PCPU_CmputNormCost_(InstSchedule *sched, COST_COMP_MODE compMode, + InstCount &execCost, bool trackCnflcts, + ParallelCPUVars &pcpu); + InstCount PCPU_CmputCost_(InstSchedule *sched, COST_COMP_MODE compMode, + InstCount &execCost, bool trackCnflcts, + ParallelCPUVars &pcpu); + void PCPU_UpdateScheduleCost(InstSchedule *sched, ParallelCPUVars &pcpu); + InstCount PCPU_CmputCostForFunction(SPILL_COST_FUNCTION SpillCF, ParallelCPUVars &pcpu); + ParallelCPUVars &GetPCPUVars(int thread); + protected: // (Chris) inline virtual const int *GetSLIL_() const { diff --git a/include/opt-sched/Scheduler/register.h b/include/opt-sched/Scheduler/register.h index 86ebf6d9..0d200f81 100644 --- a/include/opt-sched/Scheduler/register.h +++ b/include/opt-sched/Scheduler/register.h @@ -25,6 +25,10 @@ namespace opt_sched { // Forward Declaration to treat circular dependence class SchedInstruction; +struct alignas(64) AlignedInt { + int value; +}; + // Represents a a single register of a certain type and tracks the number of // times this register is defined and used. class Register { @@ -128,6 +132,13 @@ class Register { // Calls hipFree on all arrays/objects that were allocated with hipMalloc void FreeDevicePointers(); + void PCPU_AddCrntUse(int threadIdx); + void PCPU_ResetCrntUseCnt(int threadIdx); + bool PCPU_IsLive(int threadIdx) const; + void AllocParallelCPURegs(int numThreads); + void FreeParallelCPURegs(); + + private: int16_t type_; int num_; @@ -136,6 +147,9 @@ class Register { int crntUseCnt_; // Device array which holds a separate crntUseCnt_ for each thread int *dev_crntUseCnt_; + //Parallel cpu array to hold separate crntUseCnt_ + AlignedInt *pcpu_crntUseCnt_; + int crntLngth_; int physicalNumber_; BitVector conflicts_; diff --git a/lib/Scheduler/bb_spill.hip.cpp b/lib/Scheduler/bb_spill.hip.cpp index c3ed7890..05d5c559 100644 --- a/lib/Scheduler/bb_spill.hip.cpp +++ b/lib/Scheduler/bb_spill.hip.cpp @@ -557,6 +557,30 @@ void BBWithSpill::CmputCrntSpillCost_() { } #endif } + +void BBWithSpill::PCPU_CmputCrntSpillCost_(ParallelCPUVars &pcpu) { + switch (GetSpillCostFunc()) { + case SCF_PERP: + case SCF_PRP: + case SCF_PEAK_PER_TYPE: + case SCF_TARGET: + pcpu.crntSpillCost_ = pcpu.peakSpillCost_; + break; + case SCF_SUM: + pcpu.crntSpillCost_ = pcpu.totSpillCost_; + break; + case SCF_PEAK_PLUS_AVG: + pcpu.crntSpillCost_ = + pcpu.peakSpillCost_ + pcpu.totSpillCost_ / dataDepGraph_->GetInstCnt(); + break; + case SCF_SLIL: + pcpu.crntSpillCost_ = pcpu.slilSpillCost_; + break; + default: + pcpu.crntSpillCost_ = pcpu.peakSpillCost_; + break; + } +} /******************************i***********************************************/ //#define IS_DEBUG_REG_PRESSURE @@ -871,6 +895,157 @@ void BBWithSpill::UpdateSpillInfoForSchdul_(SchedInstruction *inst, schduldExitInstCnt_++; #endif } + +void BBWithSpill::PCPU_UpdateSpillInfoForSchdul_(SchedInstruction *inst, + bool trackCnflcts, + ParallelCPUVars &pcpu, + int thread) { + int16_t regType; + int defCnt, useCnt, regNum, physRegNum; + RegIndxTuple *defs, *uses; + Register *def, *use; + int liveRegs; + InstCount newSpillCost; + InstCount perpValueForSlil; + +#ifdef IS_DEBUG_REG_PRESSURE + Logger::Info("Updating reg pressure after scheduling Inst %d", + inst->GetNum()); +#endif + + defCnt = inst->GetDefs(defs); + useCnt = inst->GetUses(uses); + + // Update Live regs after uses + for (int i = 0; i < useCnt; i++) { + use = dataDepGraph_->getRegByTuple(&uses[i]); + regType = use->GetType(); + regNum = use->GetNum(); + physRegNum = use->GetPhysicalNumber(); + + if (use->IsLive() == false) { + Logger::Fatal("Reg %d of type %d is used without being defined", regNum, + regType); + } + +#ifdef IS_DEBUG_REG_PRESSURE + Logger::Info("Inst %d uses reg %d of type %d and %d uses", inst->GetNum(), + regNum, regType, use->GetUseCnt()); +#endif + + use->PCPU_AddCrntUse(thread); + + if (use->PCPU_IsLive(thread) == false) { + if (needsSLIL()) { + pcpu.sumOfLiveIntervalLengths_[regType]++; + } + + pcpu.liveRegs_[regType].SetBit(regNum, false, use->GetWght()); + +#ifdef IS_DEBUG_REG_PRESSURE + Logger::Info("Reg type %d now has %d live regs", regType, + pcpu.liveRegs_[regType].GetOneCnt()); +#endif + + if (regFiles_[regType].GetPhysRegCnt() > 0 && physRegNum >= 0) + pcpu.livePhysRegs_[regType].SetBit(physRegNum, false, use->GetWght()); + } + } + + // Update Live regs after defs + for (int i = 0; i < defCnt; i++) { + def = dataDepGraph_->getRegByTuple(&defs[i]); + regType = def->GetType(); + regNum = def->GetNum(); + physRegNum = def->GetPhysicalNumber(); + +#ifdef IS_DEBUG_REG_PRESSURE + Logger::Info("Inst %d defines reg %d of type %d and %d uses", + inst->GetNum(), regNum, regType, def->GetUseCnt()); +#endif + + if (trackCnflcts && pcpu.liveRegs_[regType].GetOneCnt() > 0) + regFiles_[regType].AddConflictsWithLiveRegs( + regNum, pcpu.liveRegs_[regType].GetOneCnt()); + + pcpu.liveRegs_[regType].SetBit(regNum, true, def->GetWght()); + +#ifdef IS_DEBUG_REG_PRESSURE + Logger::Info("Reg type %d now has %d live regs", regType, + pcpu.liveRegs_[regType].GetOneCnt()); +#endif + + if (regFiles_[regType].GetPhysRegCnt() > 0 && physRegNum >= 0) + pcpu.livePhysRegs_[regType].SetBit(physRegNum, true, def->GetWght()); + def->PCPU_ResetCrntUseCnt(thread); + } + + newSpillCost = 0; + +#ifdef IS_DEBUG_SLIL_CORRECT + if (OPTSCHED_gPrintSpills) { + Logger::Info( + "Printing live range lengths for instruction BEFORE calculation."); + for (int j = 0; j < regTypeCnt_; j++) { + Logger::Info("SLIL for regType %d %s is currently %d", j, + pcpu.sumOfLiveIntervalLengths_[j]); + } + Logger::Info("Now computing spill cost for instruction."); + } +#endif + + for (int16_t i = 0; i < regTypeCnt_; i++) { + liveRegs = pcpu.liveRegs_[i].GetWghtedCnt(); + // Set current RP for register type "i" + pcpu.regPressures_[i] = liveRegs; + // Update peak RP for register type "i" + if (liveRegs > pcpu.peakRegPressures_[i]) + pcpu.peakRegPressures_[i] = liveRegs; + + if (needsSLIL()) { + pcpu.sumOfLiveIntervalLengths_[i] += pcpu.liveRegs_[i].GetOneCnt(); + } + } + + if (GetSpillCostFunc() == SCF_SLIL) { + pcpu.slilSpillCost_ = PCPU_CmputCostForFunction(GetSpillCostFunc(), pcpu); + perpValueForSlil = PCPU_CmputCostForFunction(SCF_PERP, pcpu); + if (pcpu.peakSpillCost_ < perpValueForSlil) + pcpu.peakSpillCost_ = perpValueForSlil; + } + else + newSpillCost = PCPU_CmputCostForFunction(GetSpillCostFunc(), pcpu); + +#ifdef IS_DEBUG_SLIL_CORRECT + if (OPTSCHED_gPrintSpills) { + Logger::Info( + "Printing live range lengths for instruction AFTER calculation."); + for (int j = 0; j < regTypeCnt_; j++) { + Logger::Info("SLIL for regType %d is currently %d", j, + pcpu.sumOfLiveIntervalLengths_[j]); + } + } +#endif + + pcpu.crntStepNum_++; + pcpu.spillCosts_[pcpu.crntStepNum_] = newSpillCost; + +#ifdef IS_DEBUG_REG_PRESSURE + Logger::Info("Spill cost at step %d = %d", pcpu.crntStepNum_, newSpillCost); +#endif + + pcpu.totSpillCost_ += newSpillCost; + pcpu.peakSpillCost_ = std::max(pcpu.peakSpillCost_, newSpillCost); + + PCPU_CmputCrntSpillCost_(pcpu); + + pcpu.schduldInstCnt_++; + if (inst->MustBeInBBEntry()) + pcpu.schduldEntryInstCnt_++; + if (inst->MustBeInBBExit()) + pcpu.schduldExitInstCnt_++; +} + /*****************************************************************************/ void BBWithSpill::UpdateSpillInfoForUnSchdul_(SchedInstruction *inst) { @@ -995,6 +1170,18 @@ void BBWithSpill::SchdulInst(SchedInstruction *inst, InstCount cycleNum, UpdateSpillInfoForSchdul_(inst, trackCnflcts); } +void BBWithSpill::PCPU_SchdulInst(SchedInstruction *inst, InstCount cycleNum, + InstCount slotNum, bool trackCnflcts, + ParallelCPUVars &pcpu, + int thread) { + pcpu.crntCycleNum_ = cycleNum; + pcpu.crntSlotNum_ = slotNum; + if (inst == NULL) + return; + assert(inst != NULL); + PCPU_UpdateSpillInfoForSchdul_(inst, trackCnflcts, pcpu, thread); +} + __device__ void BBWithSpill::Dev_SchdulInst(SchedInstruction *inst, InstCount cycleNum, InstCount slotNum, bool trackCnflcts, int blockOccupancyNum) { @@ -1178,6 +1365,47 @@ InstCount BBWithSpill::CmputCostForFunction(SPILL_COST_FUNCTION SpillCF) { } } +InstCount BBWithSpill::PCPU_CmputCostForFunction(SPILL_COST_FUNCTION SpillCF, ParallelCPUVars &pcpu) { + switch (SpillCF) { + case SCF_TARGET: { + return OST->getCost(pcpu.regPressures_); + } + case SCF_SLIL: { + InstCount SLILCost = 0; + for (int i = 0; i < regTypeCnt_; i++) + SLILCost += pcpu.sumOfLiveIntervalLengths_[i]; + return SLILCost; + } + case SCF_PRP: { + InstCount PRPCost = 0; + for (int i = 0; i < regTypeCnt_; i++) + PRPCost += pcpu.regPressures_[i]; + return PRPCost; + } + case SCF_PEAK_PER_TYPE: { + InstCount SC = 0; + InstCount inc; + for (int i = 0; i < regTypeCnt_; i++) { + inc = pcpu.peakRegPressures_[i] - machMdl_->GetPhysRegCnt(i); + if (inc > 0) + SC += inc; + } + return SC; + } + default: { + // Default is PERP + InstCount inc; + InstCount SC = 0; + for (int i = 0; i < regTypeCnt_; i++) { + inc = pcpu.regPressures_[i] - machMdl_->GetPhysRegCnt(i); + if (inc > 0) + SC += inc; + } + return SC; + } + } +} + __host__ __device__ static unsigned getOccupancyWithNumVGPRs(unsigned VGPRs) { // approximation from llvm/lib/Target/AMDGPUSubtarget.cpp @@ -1750,3 +1978,143 @@ void BBWithSpill::FreeDevicePointers(int numThreads) { hipFree(dev_regPressures_); hipFree(dev_spillCosts_); } + +void BBWithSpill::AllocParallelCPUVars(int numThreads){ + pcpu_vars_ = new ParallelCPUVars[numThreads]; + for (int i = 0; i < numThreads; i++) { + // Scalar fields + pcpu_vars_[i].crntCycleNum_ = 0; + pcpu_vars_[i].crntSlotNum_ = 0; + pcpu_vars_[i].crntSpillCost_ = 0; + pcpu_vars_[i].crntStepNum_ = -1; + pcpu_vars_[i].peakSpillCost_ = 0; + pcpu_vars_[i].totSpillCost_ = 0; + pcpu_vars_[i].slilSpillCost_ = 0; + pcpu_vars_[i].dynamicSlilLowerBound_ = staticSlilLowerBound_; + pcpu_vars_[i].schduldInstCnt_ = 0; + pcpu_vars_[i].schduldEntryInstCnt_ = 0; + pcpu_vars_[i].schduldExitInstCnt_ = 0; + + // liveRegs_ — one WeightedBitVector per reg type, constructed from host liveRegs_ + pcpu_vars_[i].liveRegs_ = new WeightedBitVector[regTypeCnt_]; + for (int j = 0; j < regTypeCnt_; j++) { + pcpu_vars_[i].liveRegs_[j].Construct(regFiles_[j].GetRegCnt()); + } + + // livePhysRegs_ — only construct if physical regs exist for that type + pcpu_vars_[i].livePhysRegs_ = new WeightedBitVector[regTypeCnt_]; + for (int j = 0; j < regTypeCnt_; j++) { + int physRegCnt = regFiles_[j].GetPhysRegCnt(); + if (physRegCnt > 0) + pcpu_vars_[i].livePhysRegs_[j].Construct(physRegCnt); + } + + // peakRegPressures_ — one per reg type + pcpu_vars_[i].peakRegPressures_ = new InstCount[regTypeCnt_](); + + // regPressures_ — one per reg type + pcpu_vars_[i].regPressures_.resize(regTypeCnt_, 0); + + // spillCosts_ — one per instruction + pcpu_vars_[i].spillCosts_ = new InstCount[dataDepGraph_->GetInstCnt()](); + + // sumOfLiveIntervalLengths_ — one per reg type + pcpu_vars_[i].sumOfLiveIntervalLengths_ = new int[regTypeCnt_](); + } + for (int i = 0; i < regTypeCnt_; i++) { + for (int j = 0; j < regFiles_[i].GetRegCnt(); j++) { + regFiles_[i].GetReg(j)->AllocParallelCPURegs(numThreads); + } + } +} + +void BBWithSpill::FreeParallelCPUVars(int numThreads){ + for (int i = 0; i < numThreads; i++) { + delete[] pcpu_vars_[i].liveRegs_; + delete[] pcpu_vars_[i].livePhysRegs_; + delete[] pcpu_vars_[i].peakRegPressures_; + delete[] pcpu_vars_[i].spillCosts_; + delete[] pcpu_vars_[i].sumOfLiveIntervalLengths_; + } + delete[] pcpu_vars_; + pcpu_vars_ = nullptr; + for (int i = 0; i < regTypeCnt_; i++) { + for (int j = 0; j < regFiles_[i].GetRegCnt(); j++) { + regFiles_[i].GetReg(j)->FreeParallelCPURegs(); + } + } +} + +InstCount BBWithSpill::PCPU_GetCrntSpillCost(ParallelCPUVars &pcpu) { + return pcpu.crntSpillCost_; +} + +InstCount BBWithSpill::PCPU_ReturnPeakSpillCost(ParallelCPUVars &pcpu) { + return pcpu.peakSpillCost_; +} + +InstCount BBWithSpill::PCPU_getOccupancy(ParallelCPUVars &pcpu) { + unsigned *PRP = (unsigned *) pcpu.peakRegPressures_; + auto VGPRPressure = PRP[OptSchedDDGWrapperGCN::VGPR32]; + auto SGPRPressure = PRP[OptSchedDDGWrapperGCN::SGPR32]; + auto Occ = getAdjustedOccupancy(VGPRPressure, SGPRPressure, MaxOccLDS_, true); + return Occ; +} + +bool BBWithSpill::PCPU_closeToRPConstraint(ParallelCPUVars &pcpu, int blockOccupancyNum) { + auto Occ = getCloseToOccupancy( + pcpu.regPressures_[OptSchedDDGWrapperGCN::VGPR32], + pcpu.regPressures_[OptSchedDDGWrapperGCN::SGPR32], + MaxOccLDS_); + return Occ <= TargetOccupancy_ - blockOccupancyNum; +} + +bool BBWithSpill::PCPU_IsRPHigh(int regType, ParallelCPUVars &pcpu) const { + return pcpu.regPressures_[regType] > (unsigned int) machMdl_->GetPhysRegCnt(regType); +} + +void BBWithSpill::PCPU_UpdateScheduleCost(InstSchedule *schedule, ParallelCPUVars &pcpu) { + InstCount crntExecCost; + PCPU_CmputNormCost_(schedule, CCM_STTC, crntExecCost, false, pcpu); +} + +InstCount BBWithSpill::PCPU_CmputNormCost_(InstSchedule *sched, + COST_COMP_MODE compMode, + InstCount &execCost, bool trackCnflcts, + ParallelCPUVars &pcpu) { + InstCount cost = PCPU_CmputCost_(sched, compMode, execCost, trackCnflcts, pcpu); + + cost -= GetCostLwrBound(); + execCost -= GetExecCostLwrBound(); + + sched->SetCost(cost); + sched->SetExecCost(execCost); + sched->SetNormSpillCost(sched->GetSpillCost() * SCW_ - GetRPCostLwrBound()); + return cost; +} + +InstCount BBWithSpill::PCPU_CmputCost_(InstSchedule *sched, COST_COMP_MODE compMode, + InstCount &execCost, bool trackCnflcts, + ParallelCPUVars &pcpu) { + if (compMode == CCM_STTC) { + if (GetSpillCostFunc() == SCF_SPILLS) { + LocalRegAlloc regAlloc(sched, dataDepGraph_); + regAlloc.SetupForRegAlloc(); + regAlloc.AllocRegs(); + pcpu.crntSpillCost_ = regAlloc.GetCost(); + } + } + + assert(sched->IsComplete()); + InstCount cost = sched->GetCrntLngth() * schedCostFactor_; + execCost = cost; + cost += pcpu.crntSpillCost_ * SCW_; + sched->SetSpillCosts(pcpu.spillCosts_); + sched->SetPeakRegPressures(pcpu.peakRegPressures_); + sched->SetSpillCost(pcpu.crntSpillCost_); + return cost; +} + +ParallelCPUVars &BBWithSpill::GetPCPUVars(int thread){ + return pcpu_vars_[thread]; +} \ No newline at end of file diff --git a/lib/Scheduler/register.hip.cpp b/lib/Scheduler/register.hip.cpp index c010ab29..23441c33 100644 --- a/lib/Scheduler/register.hip.cpp +++ b/lib/Scheduler/register.hip.cpp @@ -117,6 +117,7 @@ Register &Register::operator=(const Register &rhs) { type_ = rhs.type_; useCnt_ = rhs.useCnt_; defCnt_ = rhs.defCnt_; + pcpu_crntUseCnt_ = nullptr; } return *this; @@ -196,6 +197,29 @@ Register::Register(int16_t type, int num, int physicalNumber) { isSpillCnddt_ = false; liveIn_ = false; liveOut_ = false; + + pcpu_crntUseCnt_ = nullptr; +} + +void Register::PCPU_AddCrntUse(int thread){ + pcpu_crntUseCnt_[thread].value++; +} +void Register::PCPU_ResetCrntUseCnt(int thread){ + pcpu_crntUseCnt_[thread].value = 0; +} +bool Register::PCPU_IsLive(int thread) const { + assert(pcpu_crntUseCnt_[thread].value <= useCnt_); + return pcpu_crntUseCnt_[thread].value < useCnt_; +} +void Register::AllocParallelCPURegs(int numThreads){ + pcpu_crntUseCnt_ = new AlignedInt[numThreads]; + for (int i = 0; i < numThreads; i++){ + pcpu_crntUseCnt_[i].value = 0; + } +} +void Register::FreeParallelCPURegs(){ + delete[] pcpu_crntUseCnt_; + pcpu_crntUseCnt_ = nullptr; } __host__ From 74116f559f0d23f5c37595e6eab52f198274149d Mon Sep 17 00:00:00 2001 From: Steven Inouye Date: Sun, 12 Apr 2026 13:55:47 -0700 Subject: [PATCH 03/30] Added: PCPU_FindOneSchedule --- include/opt-sched/Scheduler/aco.h | 3 + lib/Scheduler/aco.hip.cpp | 141 +++++++++++++++++++++++++++++- 2 files changed, 143 insertions(+), 1 deletion(-) diff --git a/include/opt-sched/Scheduler/aco.h b/include/opt-sched/Scheduler/aco.h index d0199de9..017f2f1d 100644 --- a/include/opt-sched/Scheduler/aco.h +++ b/include/opt-sched/Scheduler/aco.h @@ -102,6 +102,9 @@ class ACOScheduler : public ConstrainedScheduler { InstSchedule *FindManyCPUSchedule(InstCount RPTarget); + InstSchedule *PCPU_FindOneSchedule(InstCount RPTarget, + int thread, + int kernelNum = -1); __host__ __device__ diff --git a/lib/Scheduler/aco.hip.cpp b/lib/Scheduler/aco.hip.cpp index 752f2e45..e0863bb1 100644 --- a/lib/Scheduler/aco.hip.cpp +++ b/lib/Scheduler/aco.hip.cpp @@ -2430,7 +2430,7 @@ InstSchedule *ACOScheduler::FindManyCPUSchedule(InstCount RPTarget) { for (int i = 0; i < NO_CPU_THREADS; i++) { PCPUThreads.emplace_back([this, cpuScheds, i, RPTarget]() { - cpuScheds[i] = FindOneSchedule(RPTarget, NULL); + cpuScheds[i] = PCPU_FindOneSchedule(RPTarget, NULL); }); } @@ -2447,3 +2447,142 @@ InstSchedule *ACOScheduler::FindManyCPUSchedule(InstCount RPTarget) { delete[] cpuScheds; return result; } + +InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, + int thread, + int kernelNum){ + SchedInstruction *lastInst = NULL; + ACOReadyListEntry LastInstInfo; + InstSchedule *schedule; + schedule = new InstSchedule(machMdl_, dataDepGraph_, true); + bool IsSecondPass = rgn_->IsSecondPass(); + bool unnecessarilyStalling = false; + // The MaxPriority that we are getting from the ready list represents the maximum possible heuristic/key value that we can have + // I want to move all the heuristic computation stuff to another class for code tidiness reasons. + HeurType MaxPriority = kHelper1->getMaxValue(); + if (MaxPriority == 0) + MaxPriority = 1; // divide by 0 is bad + Initialize_(); + + SchedInstruction *waitFor = NULL; + InstCount waitUntil = 0; + MaxPriorityInv = 1 / (pheromone_t)MaxPriority; + + // initialize the aco ready list so that the start instruction is ready + // The luc component is 0 since the root inst uses no instructions + InstCount RootId = rootInst_->GetNum(); + HeurType RootHeuristic = kHelper1->computeKey(rootInst_, true, dataDepGraph_->RegFiles); + pheromone_t RootScore = Score(-1, RootId, RootHeuristic, !IsSecondPass); + ACOReadyListEntry InitialRoot{RootId, 0, RootHeuristic, RootScore}; + readyLs->addInstructionToReadyList(InitialRoot); + readyLs->ScoreSum = RootScore; + MaxScoringInst = 0; + lastInst = dataDepGraph_->GetInstByIndx(RootId); + bool closeToRPTarget = false; + RP0OrPositiveCount = 0; + + SchedInstruction *inst = NULL; + while (!IsSchedComplete_()) { + // incrementally calculate if there are any instructions with a neutral + // or positive effect on RP + for (InstCount I = 0; I < readyLs->getReadyListSize(); ++I) { + if (*readyLs->getInstReadyOnAtIndex(I) == crntCycleNum_) { + InstCount CandidateId = *readyLs->getInstIdAtIndex(I); + SchedInstruction *candidateInst = dataDepGraph_->GetInstByIndx(CandidateId); + HeurType candidateLUC = candidateInst->GetLastUseCnt(); + int16_t candidateDefs = candidateInst->GetDefCnt(); + if (candidateDefs <= candidateLUC) { + RP0OrPositiveCount = RP0OrPositiveCount + 1; + } + } + } + + // there are two steps to scheduling an instruction: + // 1)Select the instruction(if we are not waiting on another instruction) + inst = NULL; + if (!(waitFor && waitUntil <= crntCycleNum_)) { + // If an instruction is ready select it + assert(readyLs->getReadyListSize() > 0 || waitFor != NULL); // we should always have something in the rl + + InstCount closeToRPCheck = RPTarget - 2 < RPTarget * 9 / 10 ? RPTarget - 2 : RPTarget * 9 / 10; + closeToRPTarget = ((BBWithSpill *)rgn_)->PCPU_GetCrntSpillCost(pcpu) >= closeToRPCheck; + // select the instruction and get info on it + InstCount SelIndx = SelectInstructionC(lastInst, schedule->getTotalStalls(), rgn_, unnecessarilyStalling, closeToRPTarget, waitFor ? true: false, + readyLs, RP0OrPositiveCount, crntCycleNum_, crntSlotNum_, maxScoringInst); + + if (SelIndx != -1) { + LastInstInfo = readyLs->removeInstructionAtIndex(SelIndx); + + InstCount InstId = LastInstInfo.InstId; + inst = dataDepGraph_->GetInstByIndx(InstId); + // potentially wait on the current instruction + if (LastInstInfo.ReadyOn > crntCycleNum_ || !ChkInstLglty_(inst)) { + waitUntil = LastInstInfo.ReadyOn; + // should not wait for an instruction while already + // waiting for another instruction + assert(waitFor == NULL); + waitFor = inst; + inst = NULL; + } + + if (inst != NULL) { + #if USE_ACS + // local pheromone decay + pheromone_t *pheromone = &Pheromone(lastInst, inst, blockOccupancyNum); + *pheromone = (1 - local_decay) * *pheromone + local_decay * initialValue_; + #endif + // save the last instruction scheduled + lastInst = inst; + } + } + } + + // 2)Schedule a stall if we are still waiting, Schedule the instruction we + // are waiting for if possible, decrement waiting time + if (waitFor && waitUntil <= crntCycleNum_) { + if (ChkInstLglty_(waitFor)) { + inst = waitFor; + waitFor = NULL; + lastInst = inst; + } + } + + // boilerplate, mostly copied from ListScheduler, try not to touch it + InstCount instNum; + if (!inst) { + instNum = SCHD_STALL; + schedule->incrementTotalStalls(); + if (unnecessarilyStalling) + schedule->incrementUnnecessaryStalls(); + } else { + instNum = inst->GetNum(); + SchdulInst_(inst, crntCycleNum_); + inst->Schedule(crntCycleNum_, crntSlotNum_); + ((BBWithSpill *)rgn_)->PCPU_SchdulInst(inst, crntCycleNum_, crntSlotNum_, false, pcpu, thread); + // If an ant violates the RP cost constraint, terminate further + // schedule construction + if (((BBWithSpill*)rgn_)->PCPU_GetCrntSpillCost(pcpu) > RPTarget) { + // end schedule construction + // keep track of ants terminated + numAntsTerminated_++; + readyLs->clearReadyList(); + delete schedule; + return NULL; + } + DoRsrvSlots_(inst); + // this is annoying + UpdtSlotAvlblty_(inst); + + // new readylist update + UpdateACOReadyListC(inst, IsSecondPass, readyLs, RP0OrPositiveCount, crntCycleNum_, crntSlotNum_, maxScoringInst, 0); + } + /* Logger::Info("Chose instruction %d (for some reason)", instNum); */ + schedule->AppendInst(instNum); + if (MovToNxtSlot_(inst)) + InitNewCycle_(); + } + ((BBWithSpill *)rgn_)->PCPU_UpdateScheduleCost(schedule, pcpu); + schedule->setIsZeroPerp(((BBWithSpill *)rgn_)->PCPU_ReturnPeakSpillCost(pcpu) == 0 ); + schedule->setOccupancy(((BBWithSpill *)rgn_)->PCPU_getOccupancy(pcpu)); + return schedule; +} \ No newline at end of file From e57be448392063983356aee43c3506426e597bdc Mon Sep 17 00:00:00 2001 From: Steven Inouye Date: Sun, 12 Apr 2026 14:07:04 -0700 Subject: [PATCH 04/30] Added: 2 other PCPU functionns --- include/opt-sched/Scheduler/aco.h | 6 + lib/Scheduler/aco.hip.cpp | 196 +++++++++++++++++++++++++++++- 2 files changed, 198 insertions(+), 4 deletions(-) diff --git a/include/opt-sched/Scheduler/aco.h b/include/opt-sched/Scheduler/aco.h index 017f2f1d..af385480 100644 --- a/include/opt-sched/Scheduler/aco.h +++ b/include/opt-sched/Scheduler/aco.h @@ -105,6 +105,12 @@ class ACOScheduler : public ConstrainedScheduler { InstSchedule *PCPU_FindOneSchedule(InstCount RPTarget, int thread, int kernelNum = -1); + InstCount PCPU_SelectInstruction(SchedInstruction *lastInst, InstCount totalStalls, + SchedRegion *rgn, bool &unnecessarilyStalling, + bool closeToRPTarget, bool currentlyWaiting, + int kernelNum = -1); + inline void PCPU_UpdateACOReadyList(SchedInstruction *inst, bool IsSecondPass, int heurChoice = 0); + __host__ __device__ diff --git a/lib/Scheduler/aco.hip.cpp b/lib/Scheduler/aco.hip.cpp index e0863bb1..b3bd03d3 100644 --- a/lib/Scheduler/aco.hip.cpp +++ b/lib/Scheduler/aco.hip.cpp @@ -2507,8 +2507,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, InstCount closeToRPCheck = RPTarget - 2 < RPTarget * 9 / 10 ? RPTarget - 2 : RPTarget * 9 / 10; closeToRPTarget = ((BBWithSpill *)rgn_)->PCPU_GetCrntSpillCost(pcpu) >= closeToRPCheck; // select the instruction and get info on it - InstCount SelIndx = SelectInstructionC(lastInst, schedule->getTotalStalls(), rgn_, unnecessarilyStalling, closeToRPTarget, waitFor ? true: false, - readyLs, RP0OrPositiveCount, crntCycleNum_, crntSlotNum_, maxScoringInst); + InstCount SelIndx = PCPU_SelectInstruction(lastInst, schedule->getTotalStalls(), rgn_, unnecessarilyStalling, closeToRPTarget, waitFor ? true: false); if (SelIndx != -1) { LastInstInfo = readyLs->removeInstructionAtIndex(SelIndx); @@ -2574,7 +2573,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, UpdtSlotAvlblty_(inst); // new readylist update - UpdateACOReadyListC(inst, IsSecondPass, readyLs, RP0OrPositiveCount, crntCycleNum_, crntSlotNum_, maxScoringInst, 0); + PCPU_UpdateACOReadyList(inst, IsSecondPass, 0); } /* Logger::Info("Chose instruction %d (for some reason)", instNum); */ schedule->AppendInst(instNum); @@ -2585,4 +2584,193 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, schedule->setIsZeroPerp(((BBWithSpill *)rgn_)->PCPU_ReturnPeakSpillCost(pcpu) == 0 ); schedule->setOccupancy(((BBWithSpill *)rgn_)->PCPU_getOccupancy(pcpu)); return schedule; -} \ No newline at end of file +} +InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstCount totalStalls, + SchedRegion *rgn, bool &unnecessarilyStalling, + bool closeToRPTarget, bool currentlyWaiting, + int blockOccupancyNum){ + // if we are waiting and have no fully-ready instruction that is + // net 0 or benefit to RP, then return -1 to schedule a stall + if (currentlyWaiting && RP0OrPositiveCount == 0) + return -1; + + // calculate MaxScoringInst, and ScoreSum + pheromone_t MaxScore = -1; + InstCount MaxScoreIndx = 0; + readyLs->ScoreSum = 0; + int lastInstId = lastInst->GetNum(); + // this bool is to check if stalling could be avoided + bool couldAvoidStalling = false; + // this bool is to check if we should currently avoid unnecessary stalls + // because RP is low or we have too many stalls in the schedule + bool RPIsHigh = false; + bool tooManyStalls = totalStalls >= globalBestStalls_ * 5 / 10; + readyLs->ScoreSum = 0; + + for (InstCount I = 0; I < readyLs->getReadyListSize(); ++I) { + RPIsHigh = false; + InstCount CandidateId = *readyLs->getInstIdAtIndex(I); + SchedInstruction *candidateInst = dataDepGraph_->GetInstByIndx(CandidateId); + HeurType candidateLUC = candidateInst->GetLastUseCnt(); + int16_t candidateDefs = candidateInst->GetDefCnt(); + + // compute the score + HeurType Heur = *readyLs->getInstHeuristicAtIndex(I); + pheromone_t IScore = Score(lastInstId, *readyLs->getInstIdAtIndex(I), Heur, !rgn->IsSecondPass()); + if (RP0OrPositiveCount != 0 && candidateDefs > candidateLUC) + IScore = IScore * 9/10; + + *readyLs->getInstScoreAtIndex(I) = IScore; + readyLs->ScoreSum += IScore; + + if (currentlyWaiting) { + // if currently waiting on an instruction, do not consider semi-ready instructions + if (*readyLs->getInstReadyOnAtIndex(I) > crntCycleNum_) + continue; + + // as well as instructions with a net negative impact on RP + if (candidateDefs > candidateLUC) + continue; + } + + // add a score penalty for instructions that are not ready yet + // unnecessary stalls should not be considered if current RP is low, or if we already have too many stalls + if (*readyLs->getInstReadyOnAtIndex(I) > crntCycleNum_) { + if (RP0OrPositiveCount != 0) { + IScore = 0.0000001; + } + else { + int cyclesNeededToWait = *readyLs->getInstReadyOnAtIndex(I) - crntCycleNum_; + if (cyclesNeededToWait < globalBestStalls_) + IScore = IScore * (globalBestStalls_ - cyclesNeededToWait * 2) / globalBestStalls_; + else + IScore = IScore / globalBestStalls_; + + // check if any reg types used by the instructions are above the physical limit + SchedInstruction *tempInst = dataDepGraph_->GetInstByIndx(*readyLs->getInstIdAtIndex(I)); + RegIndxTuple *uses; + Register *use; + uint16_t usesCount = tempInst->GetUses(uses); + for (uint16_t i = 0; i < usesCount; i++) { + use = dataDepGraph_->getRegByTuple(&uses[i]); + int16_t regType = use->GetType(); + if ( ((BBWithSpill *)rgn)->IsRPHigh(regType) ) { + RPIsHigh = true; + break; + } + } + + // reduce likelihood of selecting an instruction we have to wait for IF + // RP is low or we have too many stalls already + if (!(closeToRPTarget && RPIsHigh) || tooManyStalls) { + if (globalBestStalls_ > totalStalls) + IScore = IScore * (globalBestStalls_ - totalStalls * 2) / globalBestStalls_; + else + IScore = IScore / globalBestStalls_; + } + } + } + else { + couldAvoidStalling = true; + } + + if (IScore < 0.0000001) + IScore = 0.0000001; + *readyLs->getInstScoreAtIndex(I) = IScore; + readyLs->ScoreSum += IScore; + + if(IScore > MaxScore) { + MaxScoreIndx = I; + MaxScore = IScore; + } + } + + //generate the random numbers that we will need for deciding if + //we are going to use the fixed bias or if we are going to use + //fitness proportional selection. Generate the number used for + //the fitness proportional selection point + double rand ; + pheromone_t point; + rand = RandDouble(0, 1); + point = RandDouble(0, readyLs->ScoreSum); + + //here we compute the chance that we will use fp selection or auto pick the best + double choose_best_chance; + if (use_fixed_bias) { //this is a non-diverging if stmt + choose_best_chance = (1 - (double)fixed_bias / count_) * (0 < 1 - (double)fixed_bias / count_); + } else + choose_best_chance = bias_ratio; + + //here we determine the max scoring instruction and the fp choice + //this code is a bit dense, but what we are doing is picking the + //indices of the max and fp choice instructions + //The only branch in this code is the branch for deciding to stay in the loop vs exit the loop + //this will diverge if two ants ready lists are of different sizes + // select the instruction index for fp choice + size_t fpIndx = 0; + for (size_t i = 0; i < readyLs->getReadyListSize(); ++i) { + point -= *readyLs->getInstScoreAtIndex(i); + if (point <= 0) { + fpIndx = i; + break; + } + } + + //finally we pick whether we will return the fp choice or max score inst w/o using a branch + size_t indx; + bool UseMax = (rand < choose_best_chance) || currentlyWaiting; + indx = UseMax ? MaxScoreIndx : fpIndx; + + if (couldAvoidStalling && *readyLs->getInstReadyOnAtIndex(indx) > crntCycleNum_) + unnecessarilyStalling = true; + else + unnecessarilyStalling = false; + + return indx; +} + +__host__ __device__ +inline void ACOScheduler::PCPU_UpdateACOReadyList(SchedInstruction *inst, bool IsSecondPass, int heurChoice){ + InstCount prdcsrNum, scsrRdyCycle; + // Notify each successor of this instruction that it has been scheduled. + for (SchedInstruction *crntScsr = inst->GetFrstScsr(&prdcsrNum); + crntScsr != NULL; crntScsr = inst->GetNxtScsr(&prdcsrNum)) { + bool wasLastPrdcsr = + crntScsr->PrdcsrSchduld(prdcsrNum, crntCycleNum_, scsrRdyCycle); + + if (wasLastPrdcsr) { + // If all other predecessors of this successor have been scheduled then + // we now know in which cycle this successor will become ready. + HeurType HeurWOLuc = kHelper1->computeKey(crntScsr, false, dataDepGraph_->RegFiles); + readyLs->addInstructionToReadyList(ACOReadyListEntry{crntScsr->GetNum(), scsrRdyCycle, HeurWOLuc, 0}); + } + } + + // Make sure the scores are valid. The scheduling of an instruction may + // have increased another instruction's LUC Score + PriorityEntry LUCEntry = kHelper1->getPriorityEntry(LSH_LUC); + RP0OrPositiveCount = 0; + for (InstCount I = 0; I < readyLs->getReadyListSize(); ++I) { + //we first get the heuristic without the LUC component, add the LUC + //LUC component, and then compute the score + HeurType Heur = *readyLs->getInstHeuristicAtIndex(I); + InstCount CandidateId = *readyLs->getInstIdAtIndex(I); + if (LUCEntry.Width) { + SchedInstruction *ScsrInst = dataDepGraph_->GetInstByIndx(CandidateId); + HeurType LUCVal = ScsrInst->CmputLastUseCnt(dataDepGraph_->RegFiles); + LUCVal <<= LUCEntry.Offset; + Heur &= LUCVal; + } + if (RP0OrPositiveCount) { + if (*readyLs->getInstReadyOnAtIndex(I) > crntCycleNum_) + continue; + + SchedInstruction *candidateInst = dataDepGraph_->GetInstByIndx(CandidateId); + HeurType candidateLUC = candidateInst->GetLastUseCnt(); + int16_t candidateDefs = candidateInst->GetDefCnt(); + if (candidateDefs <= candidateLUC) { + RP0OrPositiveCount = RP0OrPositiveCount + 1; + } + } + } +} From 40995278dca7a217b013eeabcd2c85c76117b730 Mon Sep 17 00:00:00 2001 From: Steven Inouye Date: Sun, 12 Apr 2026 14:15:24 -0700 Subject: [PATCH 05/30] temp --- lib/Scheduler/aco.hip.cpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/lib/Scheduler/aco.hip.cpp b/lib/Scheduler/aco.hip.cpp index b3bd03d3..3e0bf593 100644 --- a/lib/Scheduler/aco.hip.cpp +++ b/lib/Scheduler/aco.hip.cpp @@ -2427,6 +2427,7 @@ void ACOScheduler::FreeDevicePointers(bool IsSecondPass) { InstSchedule *ACOScheduler::FindManyCPUSchedule(InstCount RPTarget) { std::vector PCPUThreads; InstSchedule **cpuScheds = new InstSchedule*[NO_CPU_THREADS](); + ((BBWithSpill*)rgn_)->AllocParallelCPUVars(noCPUThreads); for (int i = 0; i < NO_CPU_THREADS; i++) { PCPUThreads.emplace_back([this, cpuScheds, i, RPTarget]() { @@ -2450,6 +2451,7 @@ InstSchedule *ACOScheduler::FindManyCPUSchedule(InstCount RPTarget) { InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, int thread, + ParallelCPUVars pcpu, int kernelNum){ SchedInstruction *lastInst = NULL; ACOReadyListEntry LastInstInfo; @@ -2462,7 +2464,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, HeurType MaxPriority = kHelper1->getMaxValue(); if (MaxPriority == 0) MaxPriority = 1; // divide by 0 is bad - Initialize_(); + Initialize_(); //For InstSchedule/ConstrainedSchedule SchedInstruction *waitFor = NULL; InstCount waitUntil = 0; From 11d759a4289db4fb77abb70d6ec5618efbd9e9d5 Mon Sep 17 00:00:00 2001 From: Steven Inouye Date: Sun, 12 Apr 2026 14:45:22 -0700 Subject: [PATCH 06/30] Tested: BBWithSpill and Register work on single thread. --- include/opt-sched/Scheduler/aco.h | 2 +- lib/Scheduler/aco.hip.cpp | 8 ++++---- 2 files changed, 5 insertions(+), 5 deletions(-) diff --git a/include/opt-sched/Scheduler/aco.h b/include/opt-sched/Scheduler/aco.h index af385480..29c38c5f 100644 --- a/include/opt-sched/Scheduler/aco.h +++ b/include/opt-sched/Scheduler/aco.h @@ -63,7 +63,6 @@ struct alignas(64) ParallelCPUVars { int schduldEntryInstCnt_; int schduldExitInstCnt_; - // These are per-thread copies of the arrays — you'll need to alloc these WeightedBitVector *liveRegs_; // array[regTypeCnt_] WeightedBitVector *livePhysRegs_; // array[regTypeCnt_] InstCount *peakRegPressures_; // array[regTypeCnt_] @@ -104,6 +103,7 @@ class ACOScheduler : public ConstrainedScheduler { InstSchedule *FindManyCPUSchedule(InstCount RPTarget); InstSchedule *PCPU_FindOneSchedule(InstCount RPTarget, int thread, + ParallelCPUVars &pcpu, int kernelNum = -1); InstCount PCPU_SelectInstruction(SchedInstruction *lastInst, InstCount totalStalls, SchedRegion *rgn, bool &unnecessarilyStalling, diff --git a/lib/Scheduler/aco.hip.cpp b/lib/Scheduler/aco.hip.cpp index 3e0bf593..e14520ae 100644 --- a/lib/Scheduler/aco.hip.cpp +++ b/lib/Scheduler/aco.hip.cpp @@ -2427,11 +2427,11 @@ void ACOScheduler::FreeDevicePointers(bool IsSecondPass) { InstSchedule *ACOScheduler::FindManyCPUSchedule(InstCount RPTarget) { std::vector PCPUThreads; InstSchedule **cpuScheds = new InstSchedule*[NO_CPU_THREADS](); - ((BBWithSpill*)rgn_)->AllocParallelCPUVars(noCPUThreads); + ((BBWithSpill*)rgn_)->AllocParallelCPUVars(NO_CPU_THREADS); for (int i = 0; i < NO_CPU_THREADS; i++) { PCPUThreads.emplace_back([this, cpuScheds, i, RPTarget]() { - cpuScheds[i] = PCPU_FindOneSchedule(RPTarget, NULL); + cpuScheds[i] = PCPU_FindOneSchedule(RPTarget, NULL, ((BBWithSpill*)rgn_)->GetPCPUVars(i)); }); } @@ -2451,7 +2451,7 @@ InstSchedule *ACOScheduler::FindManyCPUSchedule(InstCount RPTarget) { InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, int thread, - ParallelCPUVars pcpu, + ParallelCPUVars &pcpu, int kernelNum){ SchedInstruction *lastInst = NULL; ACOReadyListEntry LastInstInfo; @@ -2587,6 +2587,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, schedule->setOccupancy(((BBWithSpill *)rgn_)->PCPU_getOccupancy(pcpu)); return schedule; } + InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstCount totalStalls, SchedRegion *rgn, bool &unnecessarilyStalling, bool closeToRPTarget, bool currentlyWaiting, @@ -2731,7 +2732,6 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC return indx; } -__host__ __device__ inline void ACOScheduler::PCPU_UpdateACOReadyList(SchedInstruction *inst, bool IsSecondPass, int heurChoice){ InstCount prdcsrNum, scsrRdyCycle; // Notify each successor of this instruction that it has been scheduled. From 7caeedd2874961f42726cd05b128026442ce4112 Mon Sep 17 00:00:00 2001 From: Steven Inouye Date: Mon, 13 Apr 2026 04:18:17 -0700 Subject: [PATCH 07/30] temp --- include/opt-sched/Scheduler/aco.h | 34 ++++++++++++++++++++++++++++++- lib/Scheduler/aco.hip.cpp | 8 ++++++-- 2 files changed, 39 insertions(+), 3 deletions(-) diff --git a/include/opt-sched/Scheduler/aco.h b/include/opt-sched/Scheduler/aco.h index 29c38c5f..f460a124 100644 --- a/include/opt-sched/Scheduler/aco.h +++ b/include/opt-sched/Scheduler/aco.h @@ -60,7 +60,7 @@ struct alignas(64) ParallelCPUVars { InstCount slilSpillCost_; InstCount dynamicSlilLowerBound_; int schduldInstCnt_; - int schduldEntryInstCnt_; + int schduldEntryInstCnt_; int schduldExitInstCnt_; WeightedBitVector *liveRegs_; // array[regTypeCnt_] @@ -71,6 +71,38 @@ struct alignas(64) ParallelCPUVars { SmallVector regPressures_; // or just unsigned* }; + +struct alignas(64) PCPUACOSchedVars { + //ACOScheduler + ACOReadyList *readyLs; + InstCount maxScoringInst; + int RP0OrPositiveCount; + + //ConstrainedScheduler:InstSchedule + InstCount schduldInstCnt_; + bool isCrntCycleBlkd_; + InstCount crntCycleNum_; + InstCount crntSlotNum + ReserveSlot *rsrvSlots_; + + /* + void ConstrainedScheduler::SchdulInst_(SchedInstruction *inst, InstCount) + schduldInstCnt_ + isCrntCycleBlkd_ + void ConstrainedScheduler::DoRsrvSlots_(SchedInstruction *inst) + rsrvSlots_ + crntCycleNum_ + crntSlotNum_ + */ + + //may need multiple khelpers as well unsure +} + +struct PCPUSchedInstVars { //maybe necessary im not super confident + +} + + class ACOScheduler : public ConstrainedScheduler { public: ACOScheduler(DataDepGraph *dataDepGraph, MachineModel *machineModel, diff --git a/lib/Scheduler/aco.hip.cpp b/lib/Scheduler/aco.hip.cpp index e14520ae..c2bedcf8 100644 --- a/lib/Scheduler/aco.hip.cpp +++ b/lib/Scheduler/aco.hip.cpp @@ -2425,13 +2425,15 @@ void ACOScheduler::FreeDevicePointers(bool IsSecondPass) { } InstSchedule *ACOScheduler::FindManyCPUSchedule(InstCount RPTarget) { + Initialize_(); + std::vector PCPUThreads; InstSchedule **cpuScheds = new InstSchedule*[NO_CPU_THREADS](); ((BBWithSpill*)rgn_)->AllocParallelCPUVars(NO_CPU_THREADS); for (int i = 0; i < NO_CPU_THREADS; i++) { PCPUThreads.emplace_back([this, cpuScheds, i, RPTarget]() { - cpuScheds[i] = PCPU_FindOneSchedule(RPTarget, NULL, ((BBWithSpill*)rgn_)->GetPCPUVars(i)); + cpuScheds[i] = PCPU_FindOneSchedule(RPTarget, i, ((BBWithSpill*)rgn_)->GetPCPUVars(i)); }); } @@ -2439,6 +2441,8 @@ InstSchedule *ACOScheduler::FindManyCPUSchedule(InstCount RPTarget) { t.join(); } + ((BBWithSpill*)rgn_)->FreeParallelCPUVars(NO_CPU_THREADS); + //logic to find best schedule InstSchedule *result = cpuScheds[0]; @@ -2464,7 +2468,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, HeurType MaxPriority = kHelper1->getMaxValue(); if (MaxPriority == 0) MaxPriority = 1; // divide by 0 is bad - Initialize_(); //For InstSchedule/ConstrainedSchedule + //Initialize_(); //For InstSchedule/ConstrainedSchedule SchedInstruction *waitFor = NULL; InstCount waitUntil = 0; From 7af475cf2675565de646a7b607947520e0e23b84 Mon Sep 17 00:00:00 2001 From: Steven Inouye Date: Sat, 18 Apr 2026 14:34:42 -0700 Subject: [PATCH 08/30] Attempt 1: Allocation of ACOScheduler variables. PCPUSchedInstVars will be done at a differnt time --- include/opt-sched/Scheduler/aco.h | 23 ++++++++++++----- lib/Scheduler/aco.hip.cpp | 41 +++++++++++++++++++++++++++++++ 2 files changed, 58 insertions(+), 6 deletions(-) diff --git a/include/opt-sched/Scheduler/aco.h b/include/opt-sched/Scheduler/aco.h index f460a124..a0516f2e 100644 --- a/include/opt-sched/Scheduler/aco.h +++ b/include/opt-sched/Scheduler/aco.h @@ -75,15 +75,17 @@ struct alignas(64) ParallelCPUVars { struct alignas(64) PCPUACOSchedVars { //ACOScheduler ACOReadyList *readyLs; - InstCount maxScoringInst; + InstCount MaxScoringInst; int RP0OrPositiveCount; //ConstrainedScheduler:InstSchedule - InstCount schduldInstCnt_; - bool isCrntCycleBlkd_; - InstCount crntCycleNum_; + InstCount schduldInstCnt; + bool isCrntCycleBlkd; + InstCount crntCycleNum; InstCount crntSlotNum - ReserveSlot *rsrvSlots_; + int16_t rsrvSlotCnt; + ReserveSlot *rsrvSlots; + int16_t *avlblSlotsInCrntCycle; /* void ConstrainedScheduler::SchdulInst_(SchedInstruction *inst, InstCount) @@ -91,15 +93,21 @@ struct alignas(64) PCPUACOSchedVars { isCrntCycleBlkd_ void ConstrainedScheduler::DoRsrvSlots_(SchedInstruction *inst) rsrvSlots_ + rsrvSlotCnt_ crntCycleNum_ crntSlotNum_ + + void ConstrainedScheduler::UpdtSlotAvlblty_(SchedInstruction *inst) + avlblSlotsInCrntCycle_; */ - //may need multiple khelpers as well unsure } struct PCPUSchedInstVars { //maybe necessary im not super confident + //in FindOneSchedule + //inst->Schedule(crntCycleNum_, crntSlotNum_) + //highlights that we might need to create instructions for each thing } @@ -219,6 +227,9 @@ class ACOScheduler : public ConstrainedScheduler { KeysHelper1 *dev_kHelper1; KeysHelper2 *dev_kHelper2; InstCount *dev_MaxScoringInst; + + // ds representations for parallel CPU ACO + PCPUACOSchedVars *pcpu_sched_vars_; // True if pheromone_.elmnts_ alloced on device bool dev_pheromone_elmnts_alloced_; diff --git a/lib/Scheduler/aco.hip.cpp b/lib/Scheduler/aco.hip.cpp index c2bedcf8..e2a46940 100644 --- a/lib/Scheduler/aco.hip.cpp +++ b/lib/Scheduler/aco.hip.cpp @@ -2780,3 +2780,44 @@ inline void ACOScheduler::PCPU_UpdateACOReadyList(SchedInstruction *inst, bool I } } } + + +void ACOScheduler::AllocPCPUACOSchedVars(int numThreads){ + pcpu_sched_vars_ = new PCPUACOSchedVars[numThreads]; + for (int i = 0; i < numThreads; i++){ + pcpu_sched_vars_[i].readyLs = new ACOReadyList(dataDepGraph->GetMaxIndependentInstructions()); + pcpu_sched_vars_[i].MaxScoringInst = 0; + pcpu_sched_vars_[i].RP0OrPositiveCount = 0; + + pcpu_sched_vars_[i].schduldInstCnt = 0; + pcpu_sched_vars_[i].isCrntCycleBlkd = false; + pcpu_sched_vars_[i].crntCycleNum = 0; + pcpu_sched_vars_[i].crntSlotNum = 0; + pcpu_sched_vars_[i].rsrvSlotCnt = 0; + //i believe issuRate is safe, not looked into + pcpu_sched_vars_[i].rsrvSlots = new ReserveSlot[issuRate_]; + pcpu_sched_vars_[i].avlblSlotsInCrntCycle = new int16_t[issuTypeCnt_]; + } +} +void ACOScheduler::FreePCPUACOSchedVars(int numThreads) { + if (!pcpu_sched_vars_) + return; + + for (int i = 0; i < numThreads; i++) { + // Free ACOReadyList + delete pcpu_sched_vars_[i].readyLs; + pcpu_sched_vars_[i].readyLs = nullptr; + + // Free rsrvSlots + delete[] pcpu_sched_vars_[i].rsrvSlots; + pcpu_sched_vars_[i].rsrvSlots = nullptr; + + // Free avlblSlotsInCrntCycle + delete[] pcpu_sched_vars_[i].avlblSlotsInCrntCycle; + pcpu_sched_vars_[i].avlblSlotsInCrntCycle = nullptr; + } + + // Free the array of structs + delete[] pcpu_sched_vars_; + pcpu_sched_vars_ = nullptr; +} \ No newline at end of file From 46f52e138e63f9c8c375a3e4e6e370d66f0765d8 Mon Sep 17 00:00:00 2001 From: Steven Inouye Date: Sat, 18 Apr 2026 15:02:54 -0700 Subject: [PATCH 09/30] Attempt 1: edited allocation functions --- lib/Scheduler/aco.hip.cpp | 75 ++++++++++++++++++++++++--------------- 1 file changed, 46 insertions(+), 29 deletions(-) diff --git a/lib/Scheduler/aco.hip.cpp b/lib/Scheduler/aco.hip.cpp index e2a46940..3f898c40 100644 --- a/lib/Scheduler/aco.hip.cpp +++ b/lib/Scheduler/aco.hip.cpp @@ -2427,6 +2427,9 @@ void ACOScheduler::FreeDevicePointers(bool IsSecondPass) { InstSchedule *ACOScheduler::FindManyCPUSchedule(InstCount RPTarget) { Initialize_(); + //allocate new scheduler variables for parallel + PCPUACOSchedVars *pcpu_sched_vars = AllocPCPUACOSchedVars(NO_CPU_THREADS); + std::vector PCPUThreads; InstSchedule **cpuScheds = new InstSchedule*[NO_CPU_THREADS](); ((BBWithSpill*)rgn_)->AllocParallelCPUVars(NO_CPU_THREADS); @@ -2449,12 +2452,17 @@ InstSchedule *ACOScheduler::FindManyCPUSchedule(InstCount RPTarget) { for (int i = 1; i < NO_CPU_THREADS; i++) { delete cpuScheds[i]; } + //free new scheduler variables for parallel + FreePCPUACOSchedVars(pcpu_sched_vars, NO_CPU_THREADS); + pcpu_sched_vars = nullptr; + delete[] cpuScheds; return result; } InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, int thread, + PCPUACOSchedVars ParallelCPUVars &pcpu, int kernelNum){ SchedInstruction *lastInst = NULL; @@ -2781,43 +2789,52 @@ inline void ACOScheduler::PCPU_UpdateACOReadyList(SchedInstruction *inst, bool I } } +//allocates memory and intializes variables +PCPUACOSchedVars *ACOScheduler::AllocPCPUACOSchedVars(int numThreads) { + PCPUACOSchedVars *pcpu_sched_vars = new PCPUACOSchedVars[numThreads]; + + for (int i = 0; i < numThreads; i++) { + pcpu_sched_vars[i].readyLs = + new ACOReadyList(dataDepGraph->GetMaxIndependentInstructions()); + pcpu_sched_vars[i].MaxScoringInst = 0; + pcpu_sched_vars[i].RP0OrPositiveCount = 0; + + pcpu_sched_vars[i].schduldInstCnt = 0; + pcpu_sched_vars[i].isCrntCycleBlkd = false; + pcpu_sched_vars[i].crntCycleNum = 0; + pcpu_sched_vars[i].crntSlotNum = 0; + pcpu_sched_vars[i].rsrvSlotCnt = 0; -void ACOScheduler::AllocPCPUACOSchedVars(int numThreads){ - pcpu_sched_vars_ = new PCPUACOSchedVars[numThreads]; - for (int i = 0; i < numThreads; i++){ - pcpu_sched_vars_[i].readyLs = new ACOReadyList(dataDepGraph->GetMaxIndependentInstructions()); - pcpu_sched_vars_[i].MaxScoringInst = 0; - pcpu_sched_vars_[i].RP0OrPositiveCount = 0; - - pcpu_sched_vars_[i].schduldInstCnt = 0; - pcpu_sched_vars_[i].isCrntCycleBlkd = false; - pcpu_sched_vars_[i].crntCycleNum = 0; - pcpu_sched_vars_[i].crntSlotNum = 0; - pcpu_sched_vars_[i].rsrvSlotCnt = 0; - //i believe issuRate is safe, not looked into - pcpu_sched_vars_[i].rsrvSlots = new ReserveSlot[issuRate_]; - pcpu_sched_vars_[i].avlblSlotsInCrntCycle = new int16_t[issuTypeCnt_]; + pcpu_sched_vars[i].rsrvSlots = new ReserveSlot[issuRate_]; + pcpu_sched_vars[i].avlblSlotsInCrntCycle = new int16_t[issuTypeCnt_]; + + for (int j = 0; j < issuRate_; j++) { + pcpu_sched_vars[i].rsrvSlots[j].strtCycle = INVALID_VALUE; + pcpu_sched_vars[i].rsrvSlots[j].endCycle = INVALID_VALUE; + } + + for (int j = 0; j < issuTypeCnt_; j++) { + pcpu_sched_vars[i].avlblSlotsInCrntCycle[j] = slotsPerTypePerCycle_[j]; + } } + return pcpu_sched_vars; } -void ACOScheduler::FreePCPUACOSchedVars(int numThreads) { - if (!pcpu_sched_vars_) + +void ACOScheduler::FreePCPUACOSchedVars(PCPUACOSchedVars *pcpu_sched_vars, + int numThreads) { + if (!pcpu_sched_vars) return; for (int i = 0; i < numThreads; i++) { - // Free ACOReadyList - delete pcpu_sched_vars_[i].readyLs; - pcpu_sched_vars_[i].readyLs = nullptr; + delete pcpu_sched_vars[i].readyLs; + pcpu_sched_vars[i].readyLs = nullptr; - // Free rsrvSlots - delete[] pcpu_sched_vars_[i].rsrvSlots; - pcpu_sched_vars_[i].rsrvSlots = nullptr; + delete[] pcpu_sched_vars[i].rsrvSlots; + pcpu_sched_vars[i].rsrvSlots = nullptr; - // Free avlblSlotsInCrntCycle - delete[] pcpu_sched_vars_[i].avlblSlotsInCrntCycle; - pcpu_sched_vars_[i].avlblSlotsInCrntCycle = nullptr; + delete[] pcpu_sched_vars[i].avlblSlotsInCrntCycle; + pcpu_sched_vars[i].avlblSlotsInCrntCycle = nullptr; } - // Free the array of structs - delete[] pcpu_sched_vars_; - pcpu_sched_vars_ = nullptr; + delete[] pcpu_sched_vars; } \ No newline at end of file From fe729e90d73324a478cce72af9609f5785aa2c33 Mon Sep 17 00:00:00 2001 From: Steven Inouye Date: Sat, 18 Apr 2026 16:39:37 -0700 Subject: [PATCH 10/30] Attempt 1: Benchmarks ran initially I forgot to commit, made some changes and now its not working. --- include/opt-sched/Scheduler/aco.h | 38 ++++-- lib/Scheduler/aco.hip.cpp | 188 +++++++++++++++++++++--------- 2 files changed, 159 insertions(+), 67 deletions(-) diff --git a/include/opt-sched/Scheduler/aco.h b/include/opt-sched/Scheduler/aco.h index a0516f2e..e05096be 100644 --- a/include/opt-sched/Scheduler/aco.h +++ b/include/opt-sched/Scheduler/aco.h @@ -82,7 +82,7 @@ struct alignas(64) PCPUACOSchedVars { InstCount schduldInstCnt; bool isCrntCycleBlkd; InstCount crntCycleNum; - InstCount crntSlotNum + InstCount crntSlotNum; int16_t rsrvSlotCnt; ReserveSlot *rsrvSlots; int16_t *avlblSlotsInCrntCycle; @@ -99,16 +99,27 @@ struct alignas(64) PCPUACOSchedVars { void ConstrainedScheduler::UpdtSlotAvlblty_(SchedInstruction *inst) avlblSlotsInCrntCycle_; + + isSchedComplete ??? -> requires investigation what is totInstCnt + */ //may need multiple khelpers as well unsure -} + + /* Functions that need to be added/edited here + PCPU_FindOneSchedule + PCPU_SelectInstruction + PCPU_UpdateACOReadyList + AllocPCPUACOSchedVars + FreePCPUACOSchedVars + */ +}; struct PCPUSchedInstVars { //maybe necessary im not super confident //in FindOneSchedule //inst->Schedule(crntCycleNum_, crntSlotNum_) //highlights that we might need to create instructions for each thing -} +}; class ACOScheduler : public ConstrainedScheduler { @@ -143,15 +154,27 @@ class ACOScheduler : public ConstrainedScheduler { InstSchedule *FindManyCPUSchedule(InstCount RPTarget); InstSchedule *PCPU_FindOneSchedule(InstCount RPTarget, int thread, + PCPUACOSchedVars &pcpu_sched_vars, ParallelCPUVars &pcpu, int kernelNum = -1); InstCount PCPU_SelectInstruction(SchedInstruction *lastInst, InstCount totalStalls, SchedRegion *rgn, bool &unnecessarilyStalling, bool closeToRPTarget, bool currentlyWaiting, + int thread, + PCPUACOSchedVars &pcpu_sched_vars, int kernelNum = -1); - inline void PCPU_UpdateACOReadyList(SchedInstruction *inst, bool IsSecondPass, int heurChoice = 0); - - + inline void PCPU_UpdateACOReadyList(SchedInstruction *inst, bool IsSecondPass, + int thread, + PCPUACOSchedVars &pcpu_sched_vars, + int heurChoice = 0); + PCPUACOSchedVars *AllocPCPUACOSchedVars(int numThreads); + void FreePCPUACOSchedVars(PCPUACOSchedVars *pcpu_sched_vars, int numThreads); + /* + void PCPU_DoRsrvSlots_(SchedInstruction *inst, PCPUACOSchedVars &pcpu_sched_vars); + void PCPU_SchdulInst_(SchedInstruction *inst, PCPUACOSchedVars &pcpu_sched_vars); + void PCPU_UpdtSlotAvlblty_(SchedInstruction *inst, PCPUACOSchedVars &pcpu_sched_vars); + bool PCPU_IsSchedComplete_(PCPUACOSchedVars &pcpu_sched_vars); + */ __host__ __device__ InstSchedule *FindOneSchedule(InstCount RPTarget, @@ -227,9 +250,6 @@ class ACOScheduler : public ConstrainedScheduler { KeysHelper1 *dev_kHelper1; KeysHelper2 *dev_kHelper2; InstCount *dev_MaxScoringInst; - - // ds representations for parallel CPU ACO - PCPUACOSchedVars *pcpu_sched_vars_; // True if pheromone_.elmnts_ alloced on device bool dev_pheromone_elmnts_alloced_; diff --git a/lib/Scheduler/aco.hip.cpp b/lib/Scheduler/aco.hip.cpp index 3f898c40..9eab3124 100644 --- a/lib/Scheduler/aco.hip.cpp +++ b/lib/Scheduler/aco.hip.cpp @@ -2425,6 +2425,7 @@ void ACOScheduler::FreeDevicePointers(bool IsSecondPass) { } InstSchedule *ACOScheduler::FindManyCPUSchedule(InstCount RPTarget) { + //printf("Starting FindManyCPUSchedule"); Initialize_(); //allocate new scheduler variables for parallel @@ -2435,15 +2436,16 @@ InstSchedule *ACOScheduler::FindManyCPUSchedule(InstCount RPTarget) { ((BBWithSpill*)rgn_)->AllocParallelCPUVars(NO_CPU_THREADS); for (int i = 0; i < NO_CPU_THREADS; i++) { - PCPUThreads.emplace_back([this, cpuScheds, i, RPTarget]() { - cpuScheds[i] = PCPU_FindOneSchedule(RPTarget, i, ((BBWithSpill*)rgn_)->GetPCPUVars(i)); + PCPUThreads.emplace_back([this, cpuScheds, i, RPTarget, pcpu_sched_vars]() { + cpuScheds[i] = PCPU_FindOneSchedule(RPTarget, i, + pcpu_sched_vars[i], + ((BBWithSpill*)rgn_)->GetPCPUVars(i)); }); } for (auto &t : PCPUThreads) { t.join(); } - ((BBWithSpill*)rgn_)->FreeParallelCPUVars(NO_CPU_THREADS); //logic to find best schedule @@ -2455,16 +2457,18 @@ InstSchedule *ACOScheduler::FindManyCPUSchedule(InstCount RPTarget) { //free new scheduler variables for parallel FreePCPUACOSchedVars(pcpu_sched_vars, NO_CPU_THREADS); pcpu_sched_vars = nullptr; - delete[] cpuScheds; + + //printf("Ending FindManyCPUSchedule"); return result; } InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, int thread, - PCPUACOSchedVars + PCPUACOSchedVars &pcpu_sched_vars, ParallelCPUVars &pcpu, - int kernelNum){ + int kernelNum){ + SchedInstruction *lastInst = NULL; ACOReadyListEntry LastInstInfo; InstSchedule *schedule; @@ -2478,6 +2482,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, MaxPriority = 1; // divide by 0 is bad //Initialize_(); //For InstSchedule/ConstrainedSchedule + printf("Finding one schedule"); SchedInstruction *waitFor = NULL; InstCount waitUntil = 0; MaxPriorityInv = 1 / (pheromone_t)MaxPriority; @@ -2488,25 +2493,25 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, HeurType RootHeuristic = kHelper1->computeKey(rootInst_, true, dataDepGraph_->RegFiles); pheromone_t RootScore = Score(-1, RootId, RootHeuristic, !IsSecondPass); ACOReadyListEntry InitialRoot{RootId, 0, RootHeuristic, RootScore}; - readyLs->addInstructionToReadyList(InitialRoot); - readyLs->ScoreSum = RootScore; - MaxScoringInst = 0; + pcpu_sched_vars.readyLs->addInstructionToReadyList(InitialRoot); + pcpu_sched_vars.readyLs->ScoreSum = RootScore; + pcpu_sched_vars.MaxScoringInst = 0; lastInst = dataDepGraph_->GetInstByIndx(RootId); bool closeToRPTarget = false; - RP0OrPositiveCount = 0; + pcpu_sched_vars.RP0OrPositiveCount = 0; SchedInstruction *inst = NULL; while (!IsSchedComplete_()) { // incrementally calculate if there are any instructions with a neutral // or positive effect on RP - for (InstCount I = 0; I < readyLs->getReadyListSize(); ++I) { - if (*readyLs->getInstReadyOnAtIndex(I) == crntCycleNum_) { - InstCount CandidateId = *readyLs->getInstIdAtIndex(I); + for (InstCount I = 0; I < pcpu_sched_vars.readyLs->getReadyListSize(); ++I) { + if (*pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) == pcpu_sched_vars.crntCycleNum) { + InstCount CandidateId = *pcpu_sched_vars.readyLs->getInstIdAtIndex(I); SchedInstruction *candidateInst = dataDepGraph_->GetInstByIndx(CandidateId); HeurType candidateLUC = candidateInst->GetLastUseCnt(); int16_t candidateDefs = candidateInst->GetDefCnt(); if (candidateDefs <= candidateLUC) { - RP0OrPositiveCount = RP0OrPositiveCount + 1; + pcpu_sched_vars.RP0OrPositiveCount = pcpu_sched_vars.RP0OrPositiveCount + 1; } } } @@ -2514,22 +2519,23 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, // there are two steps to scheduling an instruction: // 1)Select the instruction(if we are not waiting on another instruction) inst = NULL; - if (!(waitFor && waitUntil <= crntCycleNum_)) { + if (!(waitFor && waitUntil <= pcpu_sched_vars.crntCycleNum)) { // If an instruction is ready select it - assert(readyLs->getReadyListSize() > 0 || waitFor != NULL); // we should always have something in the rl + assert(pcpu_sched_vars.readyLs->getReadyListSize() > 0 || waitFor != NULL); // we should always have something in the rl InstCount closeToRPCheck = RPTarget - 2 < RPTarget * 9 / 10 ? RPTarget - 2 : RPTarget * 9 / 10; closeToRPTarget = ((BBWithSpill *)rgn_)->PCPU_GetCrntSpillCost(pcpu) >= closeToRPCheck; // select the instruction and get info on it - InstCount SelIndx = PCPU_SelectInstruction(lastInst, schedule->getTotalStalls(), rgn_, unnecessarilyStalling, closeToRPTarget, waitFor ? true: false); + InstCount SelIndx = PCPU_SelectInstruction(lastInst, schedule->getTotalStalls(), rgn_, unnecessarilyStalling, closeToRPTarget, waitFor ? true: false, + thread, pcpu_sched_vars); if (SelIndx != -1) { - LastInstInfo = readyLs->removeInstructionAtIndex(SelIndx); + LastInstInfo = pcpu_sched_vars.readyLs->removeInstructionAtIndex(SelIndx); InstCount InstId = LastInstInfo.InstId; inst = dataDepGraph_->GetInstByIndx(InstId); // potentially wait on the current instruction - if (LastInstInfo.ReadyOn > crntCycleNum_ || !ChkInstLglty_(inst)) { + if (LastInstInfo.ReadyOn > pcpu_sched_vars.crntCycleNum || !ChkInstLglty_(inst)) { waitUntil = LastInstInfo.ReadyOn; // should not wait for an instruction while already // waiting for another instruction @@ -2537,7 +2543,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, waitFor = inst; inst = NULL; } - + if (inst != NULL) { #if USE_ACS // local pheromone decay @@ -2552,7 +2558,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, // 2)Schedule a stall if we are still waiting, Schedule the instruction we // are waiting for if possible, decrement waiting time - if (waitFor && waitUntil <= crntCycleNum_) { + if (waitFor && waitUntil <= pcpu_sched_vars.crntCycleNum) { if (ChkInstLglty_(waitFor)) { inst = waitFor; waitFor = NULL; @@ -2569,16 +2575,16 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, schedule->incrementUnnecessaryStalls(); } else { instNum = inst->GetNum(); - SchdulInst_(inst, crntCycleNum_); - inst->Schedule(crntCycleNum_, crntSlotNum_); - ((BBWithSpill *)rgn_)->PCPU_SchdulInst(inst, crntCycleNum_, crntSlotNum_, false, pcpu, thread); + SchdulInst_(inst, pcpu_sched_vars.crntCycleNum); + inst->Schedule(pcpu_sched_vars.crntCycleNum, pcpu_sched_vars.crntSlotNum); + ((BBWithSpill *)rgn_)->PCPU_SchdulInst(inst, pcpu_sched_vars.crntCycleNum, pcpu_sched_vars.crntSlotNum, false, pcpu, thread); // If an ant violates the RP cost constraint, terminate further // schedule construction if (((BBWithSpill*)rgn_)->PCPU_GetCrntSpillCost(pcpu) > RPTarget) { // end schedule construction // keep track of ants terminated numAntsTerminated_++; - readyLs->clearReadyList(); + pcpu_sched_vars.readyLs->clearReadyList(); delete schedule; return NULL; } @@ -2587,7 +2593,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, UpdtSlotAvlblty_(inst); // new readylist update - PCPU_UpdateACOReadyList(inst, IsSecondPass, 0); + PCPU_UpdateACOReadyList(inst, IsSecondPass, thread, pcpu_sched_vars, 0); } /* Logger::Info("Chose instruction %d (for some reason)", instNum); */ schedule->AppendInst(instNum); @@ -2603,12 +2609,14 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstCount totalStalls, SchedRegion *rgn, bool &unnecessarilyStalling, bool closeToRPTarget, bool currentlyWaiting, + int thread, + PCPUACOSchedVars &pcpu_sched_vars, int blockOccupancyNum){ // if we are waiting and have no fully-ready instruction that is // net 0 or benefit to RP, then return -1 to schedule a stall - if (currentlyWaiting && RP0OrPositiveCount == 0) + if (currentlyWaiting && pcpu_sched_vars.RP0OrPositiveCount == 0) return -1; - + printf("selectinginstruction"); // calculate MaxScoringInst, and ScoreSum pheromone_t MaxScore = -1; InstCount MaxScoreIndx = 0; @@ -2622,25 +2630,25 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC bool tooManyStalls = totalStalls >= globalBestStalls_ * 5 / 10; readyLs->ScoreSum = 0; - for (InstCount I = 0; I < readyLs->getReadyListSize(); ++I) { + for (InstCount I = 0; I < pcpu_sched_vars.readyLs->getReadyListSize(); ++I) { RPIsHigh = false; - InstCount CandidateId = *readyLs->getInstIdAtIndex(I); + InstCount CandidateId = *pcpu_sched_vars.readyLs->getInstIdAtIndex(I); SchedInstruction *candidateInst = dataDepGraph_->GetInstByIndx(CandidateId); HeurType candidateLUC = candidateInst->GetLastUseCnt(); int16_t candidateDefs = candidateInst->GetDefCnt(); // compute the score - HeurType Heur = *readyLs->getInstHeuristicAtIndex(I); - pheromone_t IScore = Score(lastInstId, *readyLs->getInstIdAtIndex(I), Heur, !rgn->IsSecondPass()); - if (RP0OrPositiveCount != 0 && candidateDefs > candidateLUC) + HeurType Heur = *pcpu_sched_vars.readyLs->getInstHeuristicAtIndex(I); + pheromone_t IScore = Score(lastInstId, *pcpu_sched_vars.readyLs->getInstIdAtIndex(I), Heur, !rgn->IsSecondPass()); + if (pcpu_sched_vars.RP0OrPositiveCount != 0 && candidateDefs > candidateLUC) IScore = IScore * 9/10; - *readyLs->getInstScoreAtIndex(I) = IScore; - readyLs->ScoreSum += IScore; + *pcpu_sched_vars.readyLs->getInstScoreAtIndex(I) = IScore; + pcpu_sched_vars.readyLs->ScoreSum += IScore; if (currentlyWaiting) { // if currently waiting on an instruction, do not consider semi-ready instructions - if (*readyLs->getInstReadyOnAtIndex(I) > crntCycleNum_) + if (*pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) > pcpu_sched_vars.crntCycleNum) continue; // as well as instructions with a net negative impact on RP @@ -2650,19 +2658,19 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC // add a score penalty for instructions that are not ready yet // unnecessary stalls should not be considered if current RP is low, or if we already have too many stalls - if (*readyLs->getInstReadyOnAtIndex(I) > crntCycleNum_) { - if (RP0OrPositiveCount != 0) { + if (*pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) > pcpu_sched_vars.crntCycleNum) { + if (pcpu_sched_vars.RP0OrPositiveCount != 0) { IScore = 0.0000001; } else { - int cyclesNeededToWait = *readyLs->getInstReadyOnAtIndex(I) - crntCycleNum_; + int cyclesNeededToWait = *pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) - pcpu_sched_vars.crntCycleNum; if (cyclesNeededToWait < globalBestStalls_) IScore = IScore * (globalBestStalls_ - cyclesNeededToWait * 2) / globalBestStalls_; else IScore = IScore / globalBestStalls_; // check if any reg types used by the instructions are above the physical limit - SchedInstruction *tempInst = dataDepGraph_->GetInstByIndx(*readyLs->getInstIdAtIndex(I)); + SchedInstruction *tempInst = dataDepGraph_->GetInstByIndx(*pcpu_sched_vars.readyLs->getInstIdAtIndex(I)); RegIndxTuple *uses; Register *use; uint16_t usesCount = tempInst->GetUses(uses); @@ -2691,8 +2699,8 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC if (IScore < 0.0000001) IScore = 0.0000001; - *readyLs->getInstScoreAtIndex(I) = IScore; - readyLs->ScoreSum += IScore; + *pcpu_sched_vars.readyLs->getInstScoreAtIndex(I) = IScore; + pcpu_sched_vars.readyLs->ScoreSum += IScore; if(IScore > MaxScore) { MaxScoreIndx = I; @@ -2707,7 +2715,7 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC double rand ; pheromone_t point; rand = RandDouble(0, 1); - point = RandDouble(0, readyLs->ScoreSum); + point = RandDouble(0, pcpu_sched_vars.readyLs->ScoreSum); //here we compute the chance that we will use fp selection or auto pick the best double choose_best_chance; @@ -2723,8 +2731,8 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC //this will diverge if two ants ready lists are of different sizes // select the instruction index for fp choice size_t fpIndx = 0; - for (size_t i = 0; i < readyLs->getReadyListSize(); ++i) { - point -= *readyLs->getInstScoreAtIndex(i); + for (size_t i = 0; i < pcpu_sched_vars.readyLs->getReadyListSize(); ++i) { + point -= *pcpu_sched_vars.readyLs->getInstScoreAtIndex(i); if (point <= 0) { fpIndx = i; break; @@ -2736,7 +2744,7 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC bool UseMax = (rand < choose_best_chance) || currentlyWaiting; indx = UseMax ? MaxScoreIndx : fpIndx; - if (couldAvoidStalling && *readyLs->getInstReadyOnAtIndex(indx) > crntCycleNum_) + if (couldAvoidStalling && *pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(indx) > pcpu_sched_vars.crntCycleNum) unnecessarilyStalling = true; else unnecessarilyStalling = false; @@ -2744,46 +2752,50 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC return indx; } -inline void ACOScheduler::PCPU_UpdateACOReadyList(SchedInstruction *inst, bool IsSecondPass, int heurChoice){ +inline void ACOScheduler::PCPU_UpdateACOReadyList(SchedInstruction *inst, bool IsSecondPass, + int thread, + PCPUACOSchedVars &pcpu_sched_vars, + int heurChoice){ InstCount prdcsrNum, scsrRdyCycle; + printf("updatingreadylist"); // Notify each successor of this instruction that it has been scheduled. for (SchedInstruction *crntScsr = inst->GetFrstScsr(&prdcsrNum); crntScsr != NULL; crntScsr = inst->GetNxtScsr(&prdcsrNum)) { bool wasLastPrdcsr = - crntScsr->PrdcsrSchduld(prdcsrNum, crntCycleNum_, scsrRdyCycle); + crntScsr->PrdcsrSchduld(prdcsrNum, pcpu_sched_vars.crntCycleNum, scsrRdyCycle); if (wasLastPrdcsr) { // If all other predecessors of this successor have been scheduled then // we now know in which cycle this successor will become ready. HeurType HeurWOLuc = kHelper1->computeKey(crntScsr, false, dataDepGraph_->RegFiles); - readyLs->addInstructionToReadyList(ACOReadyListEntry{crntScsr->GetNum(), scsrRdyCycle, HeurWOLuc, 0}); + pcpu_sched_vars.readyLs->addInstructionToReadyList(ACOReadyListEntry{crntScsr->GetNum(), scsrRdyCycle, HeurWOLuc, 0}); } } // Make sure the scores are valid. The scheduling of an instruction may // have increased another instruction's LUC Score PriorityEntry LUCEntry = kHelper1->getPriorityEntry(LSH_LUC); - RP0OrPositiveCount = 0; - for (InstCount I = 0; I < readyLs->getReadyListSize(); ++I) { + pcpu_sched_vars.RP0OrPositiveCount = 0; + for (InstCount I = 0; I < pcpu_sched_vars.readyLs->getReadyListSize(); ++I) { //we first get the heuristic without the LUC component, add the LUC //LUC component, and then compute the score - HeurType Heur = *readyLs->getInstHeuristicAtIndex(I); - InstCount CandidateId = *readyLs->getInstIdAtIndex(I); + HeurType Heur = *pcpu_sched_vars.readyLs->getInstHeuristicAtIndex(I); + InstCount CandidateId = *pcpu_sched_vars.readyLs->getInstIdAtIndex(I); if (LUCEntry.Width) { SchedInstruction *ScsrInst = dataDepGraph_->GetInstByIndx(CandidateId); HeurType LUCVal = ScsrInst->CmputLastUseCnt(dataDepGraph_->RegFiles); LUCVal <<= LUCEntry.Offset; Heur &= LUCVal; } - if (RP0OrPositiveCount) { - if (*readyLs->getInstReadyOnAtIndex(I) > crntCycleNum_) + if (pcpu_sched_vars.RP0OrPositiveCount) { + if (*pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) > pcpu_sched_vars.crntCycleNum) continue; SchedInstruction *candidateInst = dataDepGraph_->GetInstByIndx(CandidateId); HeurType candidateLUC = candidateInst->GetLastUseCnt(); int16_t candidateDefs = candidateInst->GetDefCnt(); if (candidateDefs <= candidateLUC) { - RP0OrPositiveCount = RP0OrPositiveCount + 1; + pcpu_sched_vars.RP0OrPositiveCount = pcpu_sched_vars.RP0OrPositiveCount + 1; } } } @@ -2795,7 +2807,7 @@ PCPUACOSchedVars *ACOScheduler::AllocPCPUACOSchedVars(int numThreads) { for (int i = 0; i < numThreads; i++) { pcpu_sched_vars[i].readyLs = - new ACOReadyList(dataDepGraph->GetMaxIndependentInstructions()); + new ACOReadyList(dataDepGraph_->GetMaxIndependentInstructions()); pcpu_sched_vars[i].MaxScoringInst = 0; pcpu_sched_vars[i].RP0OrPositiveCount = 0; @@ -2837,4 +2849,64 @@ void ACOScheduler::FreePCPUACOSchedVars(PCPUACOSchedVars *pcpu_sched_vars, } delete[] pcpu_sched_vars; -} \ No newline at end of file +} + +/* +void ACOScheduler::PCPU_DoRsrvSlots_(SchedInstruction *inst, PCPUACOSchedVars &pcpu_sched_vars) { + if (inst == NULL) + return; + + if (!inst->IsPipelined()) { + if (pcpu_sched_vars.rsrvSlots == NULL) + AllocRsrvSlots_(); + pcpu_sched_vars.rsrvSlots[pcpu_sched_vars.crntSlotNum].strtCycle = pcpu_sched_vars.crntCycleNum; + pcpu_sched_vars.rsrvSlots[pcpu_sched_vars.crntSlotNum].endCycle = pcpu_sched_vars.crntCycleNum + inst->GetMaxLtncy() - 1; + pcpu_sched_vars.rsrvSlotCnt++; + } +} + +void ACOScheduler::PCPU_SchdulInst_(SchedInstruction *inst, PCPUACOSchedVars &pcpu_sched_vars) { + InstCount prdcsrNum, scsrRdyCycle; + + // Notify each successor of this instruction that it has been scheduled. + if(!IsACO) { + for (SchedInstruction *crntScsr = inst->GetFrstScsr(&prdcsrNum); + crntScsr != NULL; crntScsr = inst->GetNxtScsr(&prdcsrNum)) { + bool wasLastPrdcsr = + crntScsr->PrdcsrSchduld(prdcsrNum, pcpu_sched_vars.crntCycleNum, scsrRdyCycle); + + if (wasLastPrdcsr) { + // If all other predecessors of this successor have been scheduled then + // we now know in which cycle this successor will become ready. + assert(scsrRdyCycle < schedUprBound_); + // If the first-ready list of that cycle has not been created yet. + if (frstRdyLstPerCycle_[scsrRdyCycle] == NULL) { + frstRdyLstPerCycle_[scsrRdyCycle] = + new ArrayList(dataDepGraph_->GetInstCnt()); + } + // Add this succesor to the first-ready list of the future cycle + // in which we now know it will become ready + frstRdyLstPerCycle_[scsrRdyCycle]->InsrtElmnt(crntScsr->GetNum()); + } + } + } + if (inst->BlocksCycle()) { + pcpu_sched_vars.isCrntCycleBlkd = true; + } + pcpu_sched_vars.schduldInstCnt++; +} + +void ACOScheduler::PCPU_UpdtSlotAvlblty_(SchedInstruction *inst, PCPUACOSchedVars &pcpu_sched_vars) { + if (inst == NULL) + return; + IssueType issuType = inst->GetIssueType(); + assert(issuType < issuTypeCnt_); + + assert(pcpu_sched_vars.avlblSlotsInCrntCycle[issuType] > 0); + pcpu_sched_vars.avlblSlotsInCrntCycle[issuType]--; +} + +bool ACOScheduler::PCPU_IsSchedComplete_(PCPUACOSchedVars &pcpu_sched_vars) { + return pcpu_sched_vars.schduldInstCnt == totInstCnt_; +} +*/ \ No newline at end of file From 56f8cc8dd94bb28fddb399b343500cec2397b1e5 Mon Sep 17 00:00:00 2001 From: Steven Inouye Date: Sat, 18 Apr 2026 16:54:26 -0700 Subject: [PATCH 11/30] temp --- lib/Scheduler/aco.hip.cpp | 52 +++++++++++++++++++-------------------- 1 file changed, 26 insertions(+), 26 deletions(-) diff --git a/lib/Scheduler/aco.hip.cpp b/lib/Scheduler/aco.hip.cpp index 9eab3124..36163a34 100644 --- a/lib/Scheduler/aco.hip.cpp +++ b/lib/Scheduler/aco.hip.cpp @@ -2194,7 +2194,7 @@ inline void ACOScheduler::UpdateACOReadyList(SchedInstruction *inst, bool IsSeco // Make sure the scores are valid. The scheduling of an instruction may // have increased another instruction's LUC Score PriorityEntry LUCEntry = kHelper1->getPriorityEntry(LSH_LUC); - RP0OrPositiveCount = 0; + RP0OrPositiveCount_ = 0; for (InstCount I = 0; I < readyLs->getReadyListSize(); ++I) { //we first get the heuristic without the LUC component, add the LUC //LUC component, and then compute the score @@ -2206,7 +2206,7 @@ inline void ACOScheduler::UpdateACOReadyList(SchedInstruction *inst, bool IsSeco LUCVal <<= LUCEntry.Offset; Heur &= LUCVal; } - if (RP0OrPositiveCount) { + if (RP0OrPositiveCount_) { if (*dev_readyLs->getInstReadyOnAtIndex(I) > crntCycleNum_) continue; @@ -2214,7 +2214,7 @@ inline void ACOScheduler::UpdateACOReadyList(SchedInstruction *inst, bool IsSeco HeurType candidateLUC = candidateInst->GetLastUseCnt(); int16_t candidateDefs = candidateInst->GetDefCnt(); if (candidateDefs <= candidateLUC) { - RP0OrPositiveCount = RP0OrPositiveCount + 1; + RP0OrPositiveCount_ = RP0OrPositiveCount_ + 1; } } } @@ -2425,7 +2425,7 @@ void ACOScheduler::FreeDevicePointers(bool IsSecondPass) { } InstSchedule *ACOScheduler::FindManyCPUSchedule(InstCount RPTarget) { - //printf("Starting FindManyCPUSchedule"); + printf("Starting FindManyCPUSchedule"); Initialize_(); //allocate new scheduler variables for parallel @@ -2495,23 +2495,23 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, ACOReadyListEntry InitialRoot{RootId, 0, RootHeuristic, RootScore}; pcpu_sched_vars.readyLs->addInstructionToReadyList(InitialRoot); pcpu_sched_vars.readyLs->ScoreSum = RootScore; - pcpu_sched_vars.MaxScoringInst = 0; + MaxScoringInst = 0; lastInst = dataDepGraph_->GetInstByIndx(RootId); bool closeToRPTarget = false; - pcpu_sched_vars.RP0OrPositiveCount = 0; + RP0OrPositiveCount_ = 0; SchedInstruction *inst = NULL; while (!IsSchedComplete_()) { // incrementally calculate if there are any instructions with a neutral // or positive effect on RP for (InstCount I = 0; I < pcpu_sched_vars.readyLs->getReadyListSize(); ++I) { - if (*pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) == pcpu_sched_vars.crntCycleNum) { + if (*pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) == crntCycleNum_) { InstCount CandidateId = *pcpu_sched_vars.readyLs->getInstIdAtIndex(I); SchedInstruction *candidateInst = dataDepGraph_->GetInstByIndx(CandidateId); HeurType candidateLUC = candidateInst->GetLastUseCnt(); int16_t candidateDefs = candidateInst->GetDefCnt(); if (candidateDefs <= candidateLUC) { - pcpu_sched_vars.RP0OrPositiveCount = pcpu_sched_vars.RP0OrPositiveCount + 1; + RP0OrPositiveCount_ = RP0OrPositiveCount_ + 1; } } } @@ -2519,7 +2519,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, // there are two steps to scheduling an instruction: // 1)Select the instruction(if we are not waiting on another instruction) inst = NULL; - if (!(waitFor && waitUntil <= pcpu_sched_vars.crntCycleNum)) { + if (!(waitFor && waitUntil <= crntCycleNum_)) { // If an instruction is ready select it assert(pcpu_sched_vars.readyLs->getReadyListSize() > 0 || waitFor != NULL); // we should always have something in the rl @@ -2535,7 +2535,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, InstCount InstId = LastInstInfo.InstId; inst = dataDepGraph_->GetInstByIndx(InstId); // potentially wait on the current instruction - if (LastInstInfo.ReadyOn > pcpu_sched_vars.crntCycleNum || !ChkInstLglty_(inst)) { + if (LastInstInfo.ReadyOn > crntCycleNum_ || !ChkInstLglty_(inst)) { waitUntil = LastInstInfo.ReadyOn; // should not wait for an instruction while already // waiting for another instruction @@ -2558,7 +2558,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, // 2)Schedule a stall if we are still waiting, Schedule the instruction we // are waiting for if possible, decrement waiting time - if (waitFor && waitUntil <= pcpu_sched_vars.crntCycleNum) { + if (waitFor && waitUntil <= crntCycleNum_) { if (ChkInstLglty_(waitFor)) { inst = waitFor; waitFor = NULL; @@ -2575,9 +2575,9 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, schedule->incrementUnnecessaryStalls(); } else { instNum = inst->GetNum(); - SchdulInst_(inst, pcpu_sched_vars.crntCycleNum); - inst->Schedule(pcpu_sched_vars.crntCycleNum, pcpu_sched_vars.crntSlotNum); - ((BBWithSpill *)rgn_)->PCPU_SchdulInst(inst, pcpu_sched_vars.crntCycleNum, pcpu_sched_vars.crntSlotNum, false, pcpu, thread); + SchdulInst_(inst, crntCycleNum_); + inst->Schedule(crntCycleNum_, crntSlotNum_); + ((BBWithSpill *)rgn_)->PCPU_SchdulInst(inst, crntCycleNum_, crntSlotNum_, false, pcpu, thread); // If an ant violates the RP cost constraint, terminate further // schedule construction if (((BBWithSpill*)rgn_)->PCPU_GetCrntSpillCost(pcpu) > RPTarget) { @@ -2614,7 +2614,7 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC int blockOccupancyNum){ // if we are waiting and have no fully-ready instruction that is // net 0 or benefit to RP, then return -1 to schedule a stall - if (currentlyWaiting && pcpu_sched_vars.RP0OrPositiveCount == 0) + if (currentlyWaiting && RP0OrPositiveCount_ == 0) return -1; printf("selectinginstruction"); // calculate MaxScoringInst, and ScoreSum @@ -2640,7 +2640,7 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC // compute the score HeurType Heur = *pcpu_sched_vars.readyLs->getInstHeuristicAtIndex(I); pheromone_t IScore = Score(lastInstId, *pcpu_sched_vars.readyLs->getInstIdAtIndex(I), Heur, !rgn->IsSecondPass()); - if (pcpu_sched_vars.RP0OrPositiveCount != 0 && candidateDefs > candidateLUC) + if (RP0OrPositiveCount_ != 0 && candidateDefs > candidateLUC) IScore = IScore * 9/10; *pcpu_sched_vars.readyLs->getInstScoreAtIndex(I) = IScore; @@ -2648,7 +2648,7 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC if (currentlyWaiting) { // if currently waiting on an instruction, do not consider semi-ready instructions - if (*pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) > pcpu_sched_vars.crntCycleNum) + if (*pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) > crntCycleNum_) continue; // as well as instructions with a net negative impact on RP @@ -2658,12 +2658,12 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC // add a score penalty for instructions that are not ready yet // unnecessary stalls should not be considered if current RP is low, or if we already have too many stalls - if (*pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) > pcpu_sched_vars.crntCycleNum) { - if (pcpu_sched_vars.RP0OrPositiveCount != 0) { + if (*pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) > crntCycleNum_) { + if (RP0OrPositiveCount_ != 0) { IScore = 0.0000001; } else { - int cyclesNeededToWait = *pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) - pcpu_sched_vars.crntCycleNum; + int cyclesNeededToWait = *pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) - crntCycleNum_; if (cyclesNeededToWait < globalBestStalls_) IScore = IScore * (globalBestStalls_ - cyclesNeededToWait * 2) / globalBestStalls_; else @@ -2744,7 +2744,7 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC bool UseMax = (rand < choose_best_chance) || currentlyWaiting; indx = UseMax ? MaxScoreIndx : fpIndx; - if (couldAvoidStalling && *pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(indx) > pcpu_sched_vars.crntCycleNum) + if (couldAvoidStalling && *pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(indx) > crntCycleNum_) unnecessarilyStalling = true; else unnecessarilyStalling = false; @@ -2762,7 +2762,7 @@ inline void ACOScheduler::PCPU_UpdateACOReadyList(SchedInstruction *inst, bool I for (SchedInstruction *crntScsr = inst->GetFrstScsr(&prdcsrNum); crntScsr != NULL; crntScsr = inst->GetNxtScsr(&prdcsrNum)) { bool wasLastPrdcsr = - crntScsr->PrdcsrSchduld(prdcsrNum, pcpu_sched_vars.crntCycleNum, scsrRdyCycle); + crntScsr->PrdcsrSchduld(prdcsrNum, crntCycleNum_, scsrRdyCycle); if (wasLastPrdcsr) { // If all other predecessors of this successor have been scheduled then @@ -2775,7 +2775,7 @@ inline void ACOScheduler::PCPU_UpdateACOReadyList(SchedInstruction *inst, bool I // Make sure the scores are valid. The scheduling of an instruction may // have increased another instruction's LUC Score PriorityEntry LUCEntry = kHelper1->getPriorityEntry(LSH_LUC); - pcpu_sched_vars.RP0OrPositiveCount = 0; + RP0OrPositiveCount_ = 0; for (InstCount I = 0; I < pcpu_sched_vars.readyLs->getReadyListSize(); ++I) { //we first get the heuristic without the LUC component, add the LUC //LUC component, and then compute the score @@ -2787,15 +2787,15 @@ inline void ACOScheduler::PCPU_UpdateACOReadyList(SchedInstruction *inst, bool I LUCVal <<= LUCEntry.Offset; Heur &= LUCVal; } - if (pcpu_sched_vars.RP0OrPositiveCount) { - if (*pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) > pcpu_sched_vars.crntCycleNum) + if (RP0OrPositiveCount_) { + if (*pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) > crntCycleNum_) continue; SchedInstruction *candidateInst = dataDepGraph_->GetInstByIndx(CandidateId); HeurType candidateLUC = candidateInst->GetLastUseCnt(); int16_t candidateDefs = candidateInst->GetDefCnt(); if (candidateDefs <= candidateLUC) { - pcpu_sched_vars.RP0OrPositiveCount = pcpu_sched_vars.RP0OrPositiveCount + 1; + RP0OrPositiveCount_ = RP0OrPositiveCount_ + 1; } } } From cc3f3fec47041be8ce86b5f57f76d81999baa2fc Mon Sep 17 00:00:00 2001 From: Steven Inouye Date: Sat, 18 Apr 2026 17:05:17 -0700 Subject: [PATCH 12/30] Attempt 1: benchmarks runnable, bbwith spill independent and acoreadylist independent --- lib/Scheduler/aco.hip.cpp | 29 ++++++++++++++--------------- 1 file changed, 14 insertions(+), 15 deletions(-) diff --git a/lib/Scheduler/aco.hip.cpp b/lib/Scheduler/aco.hip.cpp index 36163a34..e7ea600f 100644 --- a/lib/Scheduler/aco.hip.cpp +++ b/lib/Scheduler/aco.hip.cpp @@ -2194,7 +2194,7 @@ inline void ACOScheduler::UpdateACOReadyList(SchedInstruction *inst, bool IsSeco // Make sure the scores are valid. The scheduling of an instruction may // have increased another instruction's LUC Score PriorityEntry LUCEntry = kHelper1->getPriorityEntry(LSH_LUC); - RP0OrPositiveCount_ = 0; + RP0OrPositiveCount = 0; for (InstCount I = 0; I < readyLs->getReadyListSize(); ++I) { //we first get the heuristic without the LUC component, add the LUC //LUC component, and then compute the score @@ -2206,7 +2206,7 @@ inline void ACOScheduler::UpdateACOReadyList(SchedInstruction *inst, bool IsSeco LUCVal <<= LUCEntry.Offset; Heur &= LUCVal; } - if (RP0OrPositiveCount_) { + if (RP0OrPositiveCount) { if (*dev_readyLs->getInstReadyOnAtIndex(I) > crntCycleNum_) continue; @@ -2214,7 +2214,7 @@ inline void ACOScheduler::UpdateACOReadyList(SchedInstruction *inst, bool IsSeco HeurType candidateLUC = candidateInst->GetLastUseCnt(); int16_t candidateDefs = candidateInst->GetDefCnt(); if (candidateDefs <= candidateLUC) { - RP0OrPositiveCount_ = RP0OrPositiveCount_ + 1; + RP0OrPositiveCount = RP0OrPositiveCount + 1; } } } @@ -2425,7 +2425,6 @@ void ACOScheduler::FreeDevicePointers(bool IsSecondPass) { } InstSchedule *ACOScheduler::FindManyCPUSchedule(InstCount RPTarget) { - printf("Starting FindManyCPUSchedule"); Initialize_(); //allocate new scheduler variables for parallel @@ -2482,7 +2481,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, MaxPriority = 1; // divide by 0 is bad //Initialize_(); //For InstSchedule/ConstrainedSchedule - printf("Finding one schedule"); + SchedInstruction *waitFor = NULL; InstCount waitUntil = 0; MaxPriorityInv = 1 / (pheromone_t)MaxPriority; @@ -2498,7 +2497,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, MaxScoringInst = 0; lastInst = dataDepGraph_->GetInstByIndx(RootId); bool closeToRPTarget = false; - RP0OrPositiveCount_ = 0; + RP0OrPositiveCount = 0; SchedInstruction *inst = NULL; while (!IsSchedComplete_()) { @@ -2511,7 +2510,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, HeurType candidateLUC = candidateInst->GetLastUseCnt(); int16_t candidateDefs = candidateInst->GetDefCnt(); if (candidateDefs <= candidateLUC) { - RP0OrPositiveCount_ = RP0OrPositiveCount_ + 1; + RP0OrPositiveCount = RP0OrPositiveCount + 1; } } } @@ -2614,9 +2613,9 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC int blockOccupancyNum){ // if we are waiting and have no fully-ready instruction that is // net 0 or benefit to RP, then return -1 to schedule a stall - if (currentlyWaiting && RP0OrPositiveCount_ == 0) + if (currentlyWaiting && RP0OrPositiveCount == 0) return -1; - printf("selectinginstruction"); + // calculate MaxScoringInst, and ScoreSum pheromone_t MaxScore = -1; InstCount MaxScoreIndx = 0; @@ -2640,7 +2639,7 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC // compute the score HeurType Heur = *pcpu_sched_vars.readyLs->getInstHeuristicAtIndex(I); pheromone_t IScore = Score(lastInstId, *pcpu_sched_vars.readyLs->getInstIdAtIndex(I), Heur, !rgn->IsSecondPass()); - if (RP0OrPositiveCount_ != 0 && candidateDefs > candidateLUC) + if (RP0OrPositiveCount != 0 && candidateDefs > candidateLUC) IScore = IScore * 9/10; *pcpu_sched_vars.readyLs->getInstScoreAtIndex(I) = IScore; @@ -2659,7 +2658,7 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC // add a score penalty for instructions that are not ready yet // unnecessary stalls should not be considered if current RP is low, or if we already have too many stalls if (*pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) > crntCycleNum_) { - if (RP0OrPositiveCount_ != 0) { + if (RP0OrPositiveCount != 0) { IScore = 0.0000001; } else { @@ -2757,7 +2756,7 @@ inline void ACOScheduler::PCPU_UpdateACOReadyList(SchedInstruction *inst, bool I PCPUACOSchedVars &pcpu_sched_vars, int heurChoice){ InstCount prdcsrNum, scsrRdyCycle; - printf("updatingreadylist"); + // Notify each successor of this instruction that it has been scheduled. for (SchedInstruction *crntScsr = inst->GetFrstScsr(&prdcsrNum); crntScsr != NULL; crntScsr = inst->GetNxtScsr(&prdcsrNum)) { @@ -2775,7 +2774,7 @@ inline void ACOScheduler::PCPU_UpdateACOReadyList(SchedInstruction *inst, bool I // Make sure the scores are valid. The scheduling of an instruction may // have increased another instruction's LUC Score PriorityEntry LUCEntry = kHelper1->getPriorityEntry(LSH_LUC); - RP0OrPositiveCount_ = 0; + RP0OrPositiveCount = 0; for (InstCount I = 0; I < pcpu_sched_vars.readyLs->getReadyListSize(); ++I) { //we first get the heuristic without the LUC component, add the LUC //LUC component, and then compute the score @@ -2787,7 +2786,7 @@ inline void ACOScheduler::PCPU_UpdateACOReadyList(SchedInstruction *inst, bool I LUCVal <<= LUCEntry.Offset; Heur &= LUCVal; } - if (RP0OrPositiveCount_) { + if (RP0OrPositiveCount) { if (*pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) > crntCycleNum_) continue; @@ -2795,7 +2794,7 @@ inline void ACOScheduler::PCPU_UpdateACOReadyList(SchedInstruction *inst, bool I HeurType candidateLUC = candidateInst->GetLastUseCnt(); int16_t candidateDefs = candidateInst->GetDefCnt(); if (candidateDefs <= candidateLUC) { - RP0OrPositiveCount_ = RP0OrPositiveCount_ + 1; + RP0OrPositiveCount = RP0OrPositiveCount + 1; } } } From 3525a0ad6bcff6b811b0b8d901114d9221b8c032 Mon Sep 17 00:00:00 2001 From: Steven Inouye Date: Sat, 25 Apr 2026 13:33:05 -0700 Subject: [PATCH 13/30] Attempt: everything up to crntSlotnum in PCPUACOSchedVars does not result in any malloc errors, runtime errors nevertheless --- include/opt-sched/Scheduler/aco.h | 6 ++- lib/Scheduler/aco.hip.cpp | 75 ++++++++++++++++++++----------- 2 files changed, 53 insertions(+), 28 deletions(-) diff --git a/include/opt-sched/Scheduler/aco.h b/include/opt-sched/Scheduler/aco.h index e05096be..302e4802 100644 --- a/include/opt-sched/Scheduler/aco.h +++ b/include/opt-sched/Scheduler/aco.h @@ -102,6 +102,8 @@ struct alignas(64) PCPUACOSchedVars { isSchedComplete ??? -> requires investigation what is totInstCnt + bool ConstrainedScheduler::ChkInstLglty_(SchedInstruction *inst) const + isCrntCycleBlkd_ */ //may need multiple khelpers as well unsure @@ -169,12 +171,12 @@ class ACOScheduler : public ConstrainedScheduler { int heurChoice = 0); PCPUACOSchedVars *AllocPCPUACOSchedVars(int numThreads); void FreePCPUACOSchedVars(PCPUACOSchedVars *pcpu_sched_vars, int numThreads); - /* + void PCPU_DoRsrvSlots_(SchedInstruction *inst, PCPUACOSchedVars &pcpu_sched_vars); void PCPU_SchdulInst_(SchedInstruction *inst, PCPUACOSchedVars &pcpu_sched_vars); void PCPU_UpdtSlotAvlblty_(SchedInstruction *inst, PCPUACOSchedVars &pcpu_sched_vars); bool PCPU_IsSchedComplete_(PCPUACOSchedVars &pcpu_sched_vars); - */ + bool PCPU_ChkInstLglty_(SchedInstruction *inst, PCPUACOSchedVars &pcpu_sched_vars) const; __host__ __device__ InstSchedule *FindOneSchedule(InstCount RPTarget, diff --git a/lib/Scheduler/aco.hip.cpp b/lib/Scheduler/aco.hip.cpp index e7ea600f..9d4459a1 100644 --- a/lib/Scheduler/aco.hip.cpp +++ b/lib/Scheduler/aco.hip.cpp @@ -55,7 +55,7 @@ double RandDouble(double min, double max) { //#endif #define RUN_PCPU 1 -#define NO_CPU_THREADS 1 +#define NO_CPU_THREADS 2 ACOScheduler::ACOScheduler(DataDepGraph *dataDepGraph, MachineModel *machineModel, InstCount upperBound, @@ -2494,23 +2494,23 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, ACOReadyListEntry InitialRoot{RootId, 0, RootHeuristic, RootScore}; pcpu_sched_vars.readyLs->addInstructionToReadyList(InitialRoot); pcpu_sched_vars.readyLs->ScoreSum = RootScore; - MaxScoringInst = 0; + pcpu_sched_vars.MaxScoringInst = 0; lastInst = dataDepGraph_->GetInstByIndx(RootId); bool closeToRPTarget = false; - RP0OrPositiveCount = 0; + pcpu_sched_vars.RP0OrPositiveCount = 0; SchedInstruction *inst = NULL; while (!IsSchedComplete_()) { // incrementally calculate if there are any instructions with a neutral // or positive effect on RP for (InstCount I = 0; I < pcpu_sched_vars.readyLs->getReadyListSize(); ++I) { - if (*pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) == crntCycleNum_) { + if (*pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) == pcpu_sched_vars.crntCycleNum) { InstCount CandidateId = *pcpu_sched_vars.readyLs->getInstIdAtIndex(I); SchedInstruction *candidateInst = dataDepGraph_->GetInstByIndx(CandidateId); HeurType candidateLUC = candidateInst->GetLastUseCnt(); int16_t candidateDefs = candidateInst->GetDefCnt(); if (candidateDefs <= candidateLUC) { - RP0OrPositiveCount = RP0OrPositiveCount + 1; + pcpu_sched_vars.RP0OrPositiveCount = pcpu_sched_vars.RP0OrPositiveCount + 1; } } } @@ -2518,7 +2518,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, // there are two steps to scheduling an instruction: // 1)Select the instruction(if we are not waiting on another instruction) inst = NULL; - if (!(waitFor && waitUntil <= crntCycleNum_)) { + if (!(waitFor && waitUntil <= pcpu_sched_vars.crntCycleNum)) { // If an instruction is ready select it assert(pcpu_sched_vars.readyLs->getReadyListSize() > 0 || waitFor != NULL); // we should always have something in the rl @@ -2534,7 +2534,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, InstCount InstId = LastInstInfo.InstId; inst = dataDepGraph_->GetInstByIndx(InstId); // potentially wait on the current instruction - if (LastInstInfo.ReadyOn > crntCycleNum_ || !ChkInstLglty_(inst)) { + if (LastInstInfo.ReadyOn > pcpu_sched_vars.crntCycleNum || !ChkInstLglty_(inst)) { waitUntil = LastInstInfo.ReadyOn; // should not wait for an instruction while already // waiting for another instruction @@ -2557,7 +2557,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, // 2)Schedule a stall if we are still waiting, Schedule the instruction we // are waiting for if possible, decrement waiting time - if (waitFor && waitUntil <= crntCycleNum_) { + if (waitFor && waitUntil <= pcpu_sched_vars.crntCycleNum) { if (ChkInstLglty_(waitFor)) { inst = waitFor; waitFor = NULL; @@ -2574,9 +2574,9 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, schedule->incrementUnnecessaryStalls(); } else { instNum = inst->GetNum(); - SchdulInst_(inst, crntCycleNum_); - inst->Schedule(crntCycleNum_, crntSlotNum_); - ((BBWithSpill *)rgn_)->PCPU_SchdulInst(inst, crntCycleNum_, crntSlotNum_, false, pcpu, thread); + SchdulInst_(inst, pcpu_sched_vars.crntCycleNum); + inst->Schedule(pcpu_sched_vars.crntCycleNum, pcpu_sched_vars.crntSlotNum); + ((BBWithSpill *)rgn_)->PCPU_SchdulInst(inst, pcpu_sched_vars.crntCycleNum, pcpu_sched_vars.crntSlotNum, false, pcpu, thread); // If an ant violates the RP cost constraint, terminate further // schedule construction if (((BBWithSpill*)rgn_)->PCPU_GetCrntSpillCost(pcpu) > RPTarget) { @@ -2613,7 +2613,7 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC int blockOccupancyNum){ // if we are waiting and have no fully-ready instruction that is // net 0 or benefit to RP, then return -1 to schedule a stall - if (currentlyWaiting && RP0OrPositiveCount == 0) + if (currentlyWaiting && pcpu_sched_vars.RP0OrPositiveCount == 0) return -1; // calculate MaxScoringInst, and ScoreSum @@ -2639,7 +2639,7 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC // compute the score HeurType Heur = *pcpu_sched_vars.readyLs->getInstHeuristicAtIndex(I); pheromone_t IScore = Score(lastInstId, *pcpu_sched_vars.readyLs->getInstIdAtIndex(I), Heur, !rgn->IsSecondPass()); - if (RP0OrPositiveCount != 0 && candidateDefs > candidateLUC) + if (pcpu_sched_vars.RP0OrPositiveCount != 0 && candidateDefs > candidateLUC) IScore = IScore * 9/10; *pcpu_sched_vars.readyLs->getInstScoreAtIndex(I) = IScore; @@ -2647,7 +2647,7 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC if (currentlyWaiting) { // if currently waiting on an instruction, do not consider semi-ready instructions - if (*pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) > crntCycleNum_) + if (*pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) > pcpu_sched_vars.crntCycleNum) continue; // as well as instructions with a net negative impact on RP @@ -2657,12 +2657,12 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC // add a score penalty for instructions that are not ready yet // unnecessary stalls should not be considered if current RP is low, or if we already have too many stalls - if (*pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) > crntCycleNum_) { - if (RP0OrPositiveCount != 0) { + if (*pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) > pcpu_sched_vars.crntCycleNum) { + if (pcpu_sched_vars.RP0OrPositiveCount != 0) { IScore = 0.0000001; } else { - int cyclesNeededToWait = *pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) - crntCycleNum_; + int cyclesNeededToWait = *pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) - pcpu_sched_vars.crntCycleNum; if (cyclesNeededToWait < globalBestStalls_) IScore = IScore * (globalBestStalls_ - cyclesNeededToWait * 2) / globalBestStalls_; else @@ -2743,7 +2743,7 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC bool UseMax = (rand < choose_best_chance) || currentlyWaiting; indx = UseMax ? MaxScoreIndx : fpIndx; - if (couldAvoidStalling && *pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(indx) > crntCycleNum_) + if (couldAvoidStalling && *pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(indx) > pcpu_sched_vars.crntCycleNum) unnecessarilyStalling = true; else unnecessarilyStalling = false; @@ -2761,7 +2761,7 @@ inline void ACOScheduler::PCPU_UpdateACOReadyList(SchedInstruction *inst, bool I for (SchedInstruction *crntScsr = inst->GetFrstScsr(&prdcsrNum); crntScsr != NULL; crntScsr = inst->GetNxtScsr(&prdcsrNum)) { bool wasLastPrdcsr = - crntScsr->PrdcsrSchduld(prdcsrNum, crntCycleNum_, scsrRdyCycle); + crntScsr->PrdcsrSchduld(prdcsrNum, pcpu_sched_vars.crntCycleNum, scsrRdyCycle); if (wasLastPrdcsr) { // If all other predecessors of this successor have been scheduled then @@ -2774,7 +2774,7 @@ inline void ACOScheduler::PCPU_UpdateACOReadyList(SchedInstruction *inst, bool I // Make sure the scores are valid. The scheduling of an instruction may // have increased another instruction's LUC Score PriorityEntry LUCEntry = kHelper1->getPriorityEntry(LSH_LUC); - RP0OrPositiveCount = 0; + pcpu_sched_vars.RP0OrPositiveCount = 0; for (InstCount I = 0; I < pcpu_sched_vars.readyLs->getReadyListSize(); ++I) { //we first get the heuristic without the LUC component, add the LUC //LUC component, and then compute the score @@ -2786,15 +2786,15 @@ inline void ACOScheduler::PCPU_UpdateACOReadyList(SchedInstruction *inst, bool I LUCVal <<= LUCEntry.Offset; Heur &= LUCVal; } - if (RP0OrPositiveCount) { - if (*pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) > crntCycleNum_) + if (pcpu_sched_vars.RP0OrPositiveCount) { + if (*pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) > pcpu_sched_vars.crntCycleNum) continue; SchedInstruction *candidateInst = dataDepGraph_->GetInstByIndx(CandidateId); HeurType candidateLUC = candidateInst->GetLastUseCnt(); int16_t candidateDefs = candidateInst->GetDefCnt(); if (candidateDefs <= candidateLUC) { - RP0OrPositiveCount = RP0OrPositiveCount + 1; + pcpu_sched_vars.RP0OrPositiveCount = pcpu_sched_vars.RP0OrPositiveCount + 1; } } } @@ -2850,14 +2850,14 @@ void ACOScheduler::FreePCPUACOSchedVars(PCPUACOSchedVars *pcpu_sched_vars, delete[] pcpu_sched_vars; } -/* + void ACOScheduler::PCPU_DoRsrvSlots_(SchedInstruction *inst, PCPUACOSchedVars &pcpu_sched_vars) { if (inst == NULL) return; if (!inst->IsPipelined()) { if (pcpu_sched_vars.rsrvSlots == NULL) - AllocRsrvSlots_(); + AllocRsrvSlots_(); //this function also needs to be adjusted pcpu_sched_vars.rsrvSlots[pcpu_sched_vars.crntSlotNum].strtCycle = pcpu_sched_vars.crntCycleNum; pcpu_sched_vars.rsrvSlots[pcpu_sched_vars.crntSlotNum].endCycle = pcpu_sched_vars.crntCycleNum + inst->GetMaxLtncy() - 1; pcpu_sched_vars.rsrvSlotCnt++; @@ -2908,4 +2908,27 @@ void ACOScheduler::PCPU_UpdtSlotAvlblty_(SchedInstruction *inst, PCPUACOSchedVar bool ACOScheduler::PCPU_IsSchedComplete_(PCPUACOSchedVars &pcpu_sched_vars) { return pcpu_sched_vars.schduldInstCnt == totInstCnt_; } -*/ \ No newline at end of file + +bool ACOScheduler::PCPU_ChkInstLglty_(SchedInstruction *inst, + PCPUACOSchedVars &pcpu_sched_vars) const { + if (IsTriviallyLegal_(inst)) + return true; + // Account for instructions that block the whole cycle. + if (pcpu_sched_vars.isCrntCycleBlkd) + return false; + // Logger::Info("Cycle not blocked"); + if (inst->BlocksCycle() && pcpu_sched_vars.crntSlotNum != 0) + return false; + // Logger::Info("Does not block cycle"); + if (includesUnpipelined_ && pcpu_sched_vars.rsrvSlots && + pcpu_sched_vars.rsrvSlots[pcpu_sched_vars.crntSlotNum].strtCycle != INVALID_VALUE && + pcpu_sched_vars.crntCycleNum <= pcpu_sched_vars.rsrvSlots[pcpu_sched_vars.crntSlotNum].endCycle) { + return false; + } + + IssueType issuType = inst->GetIssueType(); + assert(issuType < issuTypeCnt_); + assert(pcpu_sched_vars.avlblSlotsInCrntCycle[issuType] >= 0); + // Logger::Info("avlblSlots = %d", avlblSlotsInCrntCycle_[issuType]); + return (pcpu_sched_vars.avlblSlotsInCrntCycle[issuType] > 0); +} \ No newline at end of file From bac8a7f18520500206e1e05e8c44c145a3ef518e Mon Sep 17 00:00:00 2001 From: Steven Inouye Date: Sat, 25 Apr 2026 13:59:49 -0700 Subject: [PATCH 14/30] Edited: Runs a few schedules, once it running the ants there is a malloc error --- lib/Scheduler/aco.hip.cpp | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/lib/Scheduler/aco.hip.cpp b/lib/Scheduler/aco.hip.cpp index 9d4459a1..a2bf771f 100644 --- a/lib/Scheduler/aco.hip.cpp +++ b/lib/Scheduler/aco.hip.cpp @@ -2574,7 +2574,8 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, schedule->incrementUnnecessaryStalls(); } else { instNum = inst->GetNum(); - SchdulInst_(inst, pcpu_sched_vars.crntCycleNum); + //PCPU_SchdulInst_(inst, pcpu_sched_vars); + SchdulInst(inst, pcpu_sched_vars.crntCycleNum) inst->Schedule(pcpu_sched_vars.crntCycleNum, pcpu_sched_vars.crntSlotNum); ((BBWithSpill *)rgn_)->PCPU_SchdulInst(inst, pcpu_sched_vars.crntCycleNum, pcpu_sched_vars.crntSlotNum, false, pcpu, thread); // If an ant violates the RP cost constraint, terminate further @@ -2587,7 +2588,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, delete schedule; return NULL; } - DoRsrvSlots_(inst); + PCPU_DoRsrvSlots_(inst, pcpu_sched_vars); // this is annoying UpdtSlotAvlblty_(inst); @@ -2819,6 +2820,7 @@ PCPUACOSchedVars *ACOScheduler::AllocPCPUACOSchedVars(int numThreads) { pcpu_sched_vars[i].rsrvSlots = new ReserveSlot[issuRate_]; pcpu_sched_vars[i].avlblSlotsInCrntCycle = new int16_t[issuTypeCnt_]; + //accounts for the ResetRsrvSlots Function for (int j = 0; j < issuRate_; j++) { pcpu_sched_vars[i].rsrvSlots[j].strtCycle = INVALID_VALUE; pcpu_sched_vars[i].rsrvSlots[j].endCycle = INVALID_VALUE; @@ -2857,7 +2859,7 @@ void ACOScheduler::PCPU_DoRsrvSlots_(SchedInstruction *inst, PCPUACOSchedVars &p if (!inst->IsPipelined()) { if (pcpu_sched_vars.rsrvSlots == NULL) - AllocRsrvSlots_(); //this function also needs to be adjusted + //AllocRsrvSlots_(); //not sure if this function is necessary pcpu_sched_vars.rsrvSlots[pcpu_sched_vars.crntSlotNum].strtCycle = pcpu_sched_vars.crntCycleNum; pcpu_sched_vars.rsrvSlots[pcpu_sched_vars.crntSlotNum].endCycle = pcpu_sched_vars.crntCycleNum + inst->GetMaxLtncy() - 1; pcpu_sched_vars.rsrvSlotCnt++; From 757db1cc7ce2fc4eb60b5f8dd21ab28e246454a4 Mon Sep 17 00:00:00 2001 From: Steven Inouye Date: Sat, 25 Apr 2026 18:31:51 -0700 Subject: [PATCH 15/30] Edited: believed to fully have adjused the ACOScheduler class now must create new adjusted data structures and function calls for the SchedInstruction clas --- include/opt-sched/Scheduler/aco.h | 23 +++++++--- lib/Scheduler/aco.hip.cpp | 75 ++++++++++++++++++++++++------- 2 files changed, 74 insertions(+), 24 deletions(-) diff --git a/include/opt-sched/Scheduler/aco.h b/include/opt-sched/Scheduler/aco.h index 302e4802..966f1869 100644 --- a/include/opt-sched/Scheduler/aco.h +++ b/include/opt-sched/Scheduler/aco.h @@ -86,6 +86,7 @@ struct alignas(64) PCPUACOSchedVars { int16_t rsrvSlotCnt; ReserveSlot *rsrvSlots; int16_t *avlblSlotsInCrntCycle; + /* void ConstrainedScheduler::SchdulInst_(SchedInstruction *inst, InstCount) @@ -104,16 +105,22 @@ struct alignas(64) PCPUACOSchedVars { bool ConstrainedScheduler::ChkInstLglty_(SchedInstruction *inst) const isCrntCycleBlkd_ + + bool ConstrainedScheduler::MovToNxtSlot_(SchedInstruction *inst) + crntCycleNum_ + crntSlotNum_ + crntRealSlotNum_ + + void ConstrainedScheduler::InitNewCycle_() + isCrntCycleBlkd_ + avlblSlotsInCrntCycle_ */ + + + + //may need multiple khelpers as well unsure - /* Functions that need to be added/edited here - PCPU_FindOneSchedule - PCPU_SelectInstruction - PCPU_UpdateACOReadyList - AllocPCPUACOSchedVars - FreePCPUACOSchedVars - */ }; struct PCPUSchedInstVars { //maybe necessary im not super confident @@ -177,6 +184,8 @@ class ACOScheduler : public ConstrainedScheduler { void PCPU_UpdtSlotAvlblty_(SchedInstruction *inst, PCPUACOSchedVars &pcpu_sched_vars); bool PCPU_IsSchedComplete_(PCPUACOSchedVars &pcpu_sched_vars); bool PCPU_ChkInstLglty_(SchedInstruction *inst, PCPUACOSchedVars &pcpu_sched_vars) const; + bool PCPU_MovToNxtSlot_(SchedInstruction *inst, PCPUACOSchedVars &pcpu_sched_vars); + void PCPU_InitNewCycle_(PCPUACOSchedVars &pcpu_sched_vars); __host__ __device__ InstSchedule *FindOneSchedule(InstCount RPTarget, diff --git a/lib/Scheduler/aco.hip.cpp b/lib/Scheduler/aco.hip.cpp index a2bf771f..d9ea71d5 100644 --- a/lib/Scheduler/aco.hip.cpp +++ b/lib/Scheduler/aco.hip.cpp @@ -2426,7 +2426,7 @@ void ACOScheduler::FreeDevicePointers(bool IsSecondPass) { InstSchedule *ACOScheduler::FindManyCPUSchedule(InstCount RPTarget) { Initialize_(); - + Logger::Info("FindManyCPUSchedule"); //allocate new scheduler variables for parallel PCPUACOSchedVars *pcpu_sched_vars = AllocPCPUACOSchedVars(NO_CPU_THREADS); @@ -2467,7 +2467,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, PCPUACOSchedVars &pcpu_sched_vars, ParallelCPUVars &pcpu, int kernelNum){ - + Logger::Info("FindOneSchedule, %d", thread); SchedInstruction *lastInst = NULL; ACOReadyListEntry LastInstInfo; InstSchedule *schedule; @@ -2481,13 +2481,13 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, MaxPriority = 1; // divide by 0 is bad //Initialize_(); //For InstSchedule/ConstrainedSchedule - SchedInstruction *waitFor = NULL; InstCount waitUntil = 0; MaxPriorityInv = 1 / (pheromone_t)MaxPriority; // initialize the aco ready list so that the start instruction is ready // The luc component is 0 since the root inst uses no instructions + Logger::Info("ReadyListFunctions 1, %d", thread); InstCount RootId = rootInst_->GetNum(); HeurType RootHeuristic = kHelper1->computeKey(rootInst_, true, dataDepGraph_->RegFiles); pheromone_t RootScore = Score(-1, RootId, RootHeuristic, !IsSecondPass); @@ -2499,10 +2499,13 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, bool closeToRPTarget = false; pcpu_sched_vars.RP0OrPositiveCount = 0; + Logger::Info("While not PCPU_IsSchedComplete, %d", thread); SchedInstruction *inst = NULL; - while (!IsSchedComplete_()) { + while (!PCPU_IsSchedComplete_(pcpu_sched_vars)) { // incrementally calculate if there are any instructions with a neutral // or positive effect on RP + + Logger::Info("Neut/pos on RP, %d", thread); for (InstCount I = 0; I < pcpu_sched_vars.readyLs->getReadyListSize(); ++I) { if (*pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) == pcpu_sched_vars.crntCycleNum) { InstCount CandidateId = *pcpu_sched_vars.readyLs->getInstIdAtIndex(I); @@ -2517,6 +2520,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, // there are two steps to scheduling an instruction: // 1)Select the instruction(if we are not waiting on another instruction) + Logger::Info("SelectInstruction, %d", thread); inst = NULL; if (!(waitFor && waitUntil <= pcpu_sched_vars.crntCycleNum)) { // If an instruction is ready select it @@ -2534,7 +2538,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, InstCount InstId = LastInstInfo.InstId; inst = dataDepGraph_->GetInstByIndx(InstId); // potentially wait on the current instruction - if (LastInstInfo.ReadyOn > pcpu_sched_vars.crntCycleNum || !ChkInstLglty_(inst)) { + if (LastInstInfo.ReadyOn > pcpu_sched_vars.crntCycleNum || !PCPU_ChkInstLglty_(inst, pcpu_sched_vars)) { waitUntil = LastInstInfo.ReadyOn; // should not wait for an instruction while already // waiting for another instruction @@ -2555,10 +2559,11 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, } } + Logger::Info("Scheduling Stalls, %d", thread); // 2)Schedule a stall if we are still waiting, Schedule the instruction we // are waiting for if possible, decrement waiting time if (waitFor && waitUntil <= pcpu_sched_vars.crntCycleNum) { - if (ChkInstLglty_(waitFor)) { + if (PCPU_ChkInstLglty_(waitFor, pcpu_sched_vars)) { inst = waitFor; waitFor = NULL; lastInst = inst; @@ -2566,6 +2571,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, } // boilerplate, mostly copied from ListScheduler, try not to touch it + Logger::Info("BoilerPlate, %d", thread); InstCount instNum; if (!inst) { instNum = SCHD_STALL; @@ -2574,8 +2580,8 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, schedule->incrementUnnecessaryStalls(); } else { instNum = inst->GetNum(); - //PCPU_SchdulInst_(inst, pcpu_sched_vars); - SchdulInst(inst, pcpu_sched_vars.crntCycleNum) + PCPU_SchdulInst_(inst, pcpu_sched_vars); + //SchdulInst(inst, pcpu_sched_vars.crntCycleNum) inst->Schedule(pcpu_sched_vars.crntCycleNum, pcpu_sched_vars.crntSlotNum); ((BBWithSpill *)rgn_)->PCPU_SchdulInst(inst, pcpu_sched_vars.crntCycleNum, pcpu_sched_vars.crntSlotNum, false, pcpu, thread); // If an ant violates the RP cost constraint, terminate further @@ -2590,19 +2596,22 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, } PCPU_DoRsrvSlots_(inst, pcpu_sched_vars); // this is annoying - UpdtSlotAvlblty_(inst); + PCPU_UpdtSlotAvlblty_(inst, pcpu_sched_vars); // new readylist update PCPU_UpdateACOReadyList(inst, IsSecondPass, thread, pcpu_sched_vars, 0); } /* Logger::Info("Chose instruction %d (for some reason)", instNum); */ schedule->AppendInst(instNum); - if (MovToNxtSlot_(inst)) - InitNewCycle_(); + if (PCPU_MovToNxtSlot_(inst, pcpu_sched_vars)) + PCPU_InitNewCycle_(pcpu_sched_vars); } + + Logger::Info("Updating schedule cost, %d", thread); ((BBWithSpill *)rgn_)->PCPU_UpdateScheduleCost(schedule, pcpu); schedule->setIsZeroPerp(((BBWithSpill *)rgn_)->PCPU_ReturnPeakSpillCost(pcpu) == 0 ); schedule->setOccupancy(((BBWithSpill *)rgn_)->PCPU_getOccupancy(pcpu)); + Logger::Info("Done, %d", thread); return schedule; } @@ -2620,7 +2629,7 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC // calculate MaxScoringInst, and ScoreSum pheromone_t MaxScore = -1; InstCount MaxScoreIndx = 0; - readyLs->ScoreSum = 0; + pcpu_sched_vars.readyLs->ScoreSum = 0; int lastInstId = lastInst->GetNum(); // this bool is to check if stalling could be avoided bool couldAvoidStalling = false; @@ -2628,7 +2637,7 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC // because RP is low or we have too many stalls in the schedule bool RPIsHigh = false; bool tooManyStalls = totalStalls >= globalBestStalls_ * 5 / 10; - readyLs->ScoreSum = 0; + pcpu_sched_vars.readyLs->ScoreSum = 0; for (InstCount I = 0; I < pcpu_sched_vars.readyLs->getReadyListSize(); ++I) { RPIsHigh = false; @@ -2757,7 +2766,7 @@ inline void ACOScheduler::PCPU_UpdateACOReadyList(SchedInstruction *inst, bool I PCPUACOSchedVars &pcpu_sched_vars, int heurChoice){ InstCount prdcsrNum, scsrRdyCycle; - + Logger::Info("Start UpdateACOReadyList, %d", thread); // Notify each successor of this instruction that it has been scheduled. for (SchedInstruction *crntScsr = inst->GetFrstScsr(&prdcsrNum); crntScsr != NULL; crntScsr = inst->GetNxtScsr(&prdcsrNum)) { @@ -2799,6 +2808,7 @@ inline void ACOScheduler::PCPU_UpdateACOReadyList(SchedInstruction *inst, bool I } } } + Logger::Info("Finish UpdateACOReadyList, %d", thread); } //allocates memory and intializes variables @@ -2858,14 +2868,16 @@ void ACOScheduler::PCPU_DoRsrvSlots_(SchedInstruction *inst, PCPUACOSchedVars &p return; if (!inst->IsPipelined()) { - if (pcpu_sched_vars.rsrvSlots == NULL) - //AllocRsrvSlots_(); //not sure if this function is necessary + if (pcpu_sched_vars.rsrvSlots == NULL){ + //AllocRsrvSlots_(); //not sure if this function is necessary + } pcpu_sched_vars.rsrvSlots[pcpu_sched_vars.crntSlotNum].strtCycle = pcpu_sched_vars.crntCycleNum; pcpu_sched_vars.rsrvSlots[pcpu_sched_vars.crntSlotNum].endCycle = pcpu_sched_vars.crntCycleNum + inst->GetMaxLtncy() - 1; pcpu_sched_vars.rsrvSlotCnt++; } } +/*unsur if i need to make multiple frstRdyLstPerCycle*/ void ACOScheduler::PCPU_SchdulInst_(SchedInstruction *inst, PCPUACOSchedVars &pcpu_sched_vars) { InstCount prdcsrNum, scsrRdyCycle; @@ -2882,6 +2894,7 @@ void ACOScheduler::PCPU_SchdulInst_(SchedInstruction *inst, PCPUACOSchedVars &pc assert(scsrRdyCycle < schedUprBound_); // If the first-ready list of that cycle has not been created yet. if (frstRdyLstPerCycle_[scsrRdyCycle] == NULL) { + frstRdyLstPerCycle_[scsrRdyCycle] = new ArrayList(dataDepGraph_->GetInstCnt()); } @@ -2912,7 +2925,7 @@ bool ACOScheduler::PCPU_IsSchedComplete_(PCPUACOSchedVars &pcpu_sched_vars) { } bool ACOScheduler::PCPU_ChkInstLglty_(SchedInstruction *inst, - PCPUACOSchedVars &pcpu_sched_vars) const { + PCPUACOSchedVars &pcpu_sched_vars) const { if (IsTriviallyLegal_(inst)) return true; // Account for instructions that block the whole cycle. @@ -2933,4 +2946,32 @@ bool ACOScheduler::PCPU_ChkInstLglty_(SchedInstruction *inst, assert(pcpu_sched_vars.avlblSlotsInCrntCycle[issuType] >= 0); // Logger::Info("avlblSlots = %d", avlblSlotsInCrntCycle_[issuType]); return (pcpu_sched_vars.avlblSlotsInCrntCycle[issuType] > 0); +} + + +bool ACOScheduler::PCPU_MovToNxtSlot_(SchedInstruction *inst, + PCPUACOSchedVars &pcpu_sched_vars) { + // If we are currently in the last slot of the current cycle. + // I don' tthink crntRealSlotNum_ is necessary, it is never used + if (pcpu_sched_vars.crntSlotNum == (issuRate_ - 1)) { + pcpu_sched_vars.crntCycleNum++; + pcpu_sched_vars.crntSlotNum = 0; + //crntRealSlotNum_ = 0; + return true; + } else { + pcpu_sched_vars.crntSlotNum++; + if (inst && machMdl_->IsRealInst(inst->GetInstType())) + //crntRealSlotNum_++; + return false; + } +} + + +void ACOScheduler::PCPU_InitNewCycle_(PCPUACOSchedVars &pcpu_sched_vars) { + //assert(pcpu_sched_vars.crntSlotNum == 0 && crntRealSlotNum_ == 0); + assert(pcpu_sched_vars.crntSlotNum == 0); + for (int i = 0; i < issuTypeCnt_; i++) { + pcpu_sched_vars.avlblSlotsInCrntCycle[i] = slotsPerTypePerCycle_[i]; + } + pcpu_sched_vars.isCrntCycleBlkd = false; } \ No newline at end of file From eb2554b1bfd03aa5246f7f773c983918cb573f3b Mon Sep 17 00:00:00 2001 From: Steven Inouye Date: Sat, 25 Apr 2026 21:03:46 -0700 Subject: [PATCH 16/30] Added: Data structure for SchedInstruction class --- include/opt-sched/Scheduler/aco.h | 12 ++--- .../opt-sched/Scheduler/sched_basic_data.h | 28 +++++++++++ lib/Scheduler/aco.hip.cpp | 19 +++++++- lib/Scheduler/sched_basic_data.hip.cpp | 46 +++++++++++++++++++ 4 files changed, 95 insertions(+), 10 deletions(-) diff --git a/include/opt-sched/Scheduler/aco.h b/include/opt-sched/Scheduler/aco.h index 966f1869..e8da2e51 100644 --- a/include/opt-sched/Scheduler/aco.h +++ b/include/opt-sched/Scheduler/aco.h @@ -123,14 +123,6 @@ struct alignas(64) PCPUACOSchedVars { }; -struct PCPUSchedInstVars { //maybe necessary im not super confident - - //in FindOneSchedule - //inst->Schedule(crntCycleNum_, crntSlotNum_) - //highlights that we might need to create instructions for each thing -}; - - class ACOScheduler : public ConstrainedScheduler { public: ACOScheduler(DataDepGraph *dataDepGraph, MachineModel *machineModel, @@ -186,7 +178,9 @@ class ACOScheduler : public ConstrainedScheduler { bool PCPU_ChkInstLglty_(SchedInstruction *inst, PCPUACOSchedVars &pcpu_sched_vars) const; bool PCPU_MovToNxtSlot_(SchedInstruction *inst, PCPUACOSchedVars &pcpu_sched_vars); void PCPU_InitNewCycle_(PCPUACOSchedVars &pcpu_sched_vars); - + void PCPU_InitSchedInsts(int numThreads); + void PCPU_FreeSchedInsts(int numThreads); + __host__ __device__ InstSchedule *FindOneSchedule(InstCount RPTarget, InstSchedule *dev_schedule = NULL, int kernelNum = -1); diff --git a/include/opt-sched/Scheduler/sched_basic_data.h b/include/opt-sched/Scheduler/sched_basic_data.h index 859f6d3d..e88476da 100644 --- a/include/opt-sched/Scheduler/sched_basic_data.h +++ b/include/opt-sched/Scheduler/sched_basic_data.h @@ -101,6 +101,25 @@ struct RegIndxTuple { : regType_(regType), regNum_(regNum) {} }; +//New Structure for parallel cpu +struct alignas(64) PCPUSchedInstVars { + //InitForSchduling -- where to check on how to initialize variables + //SchedInstruction InitForSchdulng + InstCount crntSchedCycle; + int16_t lastUseCnt; + bool ready; + InstCount minRdyCycle; + InstCount unschduldPrdcsrCnt; + InstCount unschduldScsrCnt; + InstCount crntRlxdCycle; + InstCount *rdyCyclePerPrdcsr; + InstCount *prevMinRdyCyclePerPrdcsr; + + //not initialized in that thing for some reason + InstCount crntSchedSlot; + +}; + // The type of instruction signatures, used by the enumerator's history table to // keep track of partial schedules. typedef UDT_HASHKEY InstSignature; @@ -558,6 +577,12 @@ class SchedInstruction : public GraphNode { // parallel ACO on device void AllocDevArraysForParallelACO(int numThreads); + + void AllocPCPUVars(int numThreads); + void InitPCPUVars(int numThreads); + void FreePCPUVars(int numThreds); + + friend class SchedRange; // This instruction's index in the scsrs_, latencies_, predOrder_ arrays @@ -741,6 +766,9 @@ class SchedInstruction : public GraphNode { // much faster copying to the Device SchedInstruction *insts_; + //For parallel CPU + PCPUSchedInstVars *pcpu_inst_vars_; + // TODO(ghassan): Document. __host__ InstCount CmputCrtclPath_(DIRECTION dir, SchedInstruction *ref = NULL); diff --git a/lib/Scheduler/aco.hip.cpp b/lib/Scheduler/aco.hip.cpp index d9ea71d5..d6ca0b7d 100644 --- a/lib/Scheduler/aco.hip.cpp +++ b/lib/Scheduler/aco.hip.cpp @@ -2426,6 +2426,7 @@ void ACOScheduler::FreeDevicePointers(bool IsSecondPass) { InstSchedule *ACOScheduler::FindManyCPUSchedule(InstCount RPTarget) { Initialize_(); + PCPU_InitSchedInsts(NO_CPU_THREADS); Logger::Info("FindManyCPUSchedule"); //allocate new scheduler variables for parallel PCPUACOSchedVars *pcpu_sched_vars = AllocPCPUACOSchedVars(NO_CPU_THREADS); @@ -2457,6 +2458,7 @@ InstSchedule *ACOScheduler::FindManyCPUSchedule(InstCount RPTarget) { FreePCPUACOSchedVars(pcpu_sched_vars, NO_CPU_THREADS); pcpu_sched_vars = nullptr; delete[] cpuScheds; + PCPU_FreeSchedInsts(NO_CPU_THREADS); //printf("Ending FindManyCPUSchedule"); return result; @@ -2881,7 +2883,8 @@ void ACOScheduler::PCPU_DoRsrvSlots_(SchedInstruction *inst, PCPUACOSchedVars &p void ACOScheduler::PCPU_SchdulInst_(SchedInstruction *inst, PCPUACOSchedVars &pcpu_sched_vars) { InstCount prdcsrNum, scsrRdyCycle; - // Notify each successor of this instruction that it has been scheduled. + // Notify each successor of this instruction that it has been scheduled.' + // if not running aco if(!IsACO) { for (SchedInstruction *crntScsr = inst->GetFrstScsr(&prdcsrNum); crntScsr != NULL; crntScsr = inst->GetNxtScsr(&prdcsrNum)) { @@ -2974,4 +2977,18 @@ void ACOScheduler::PCPU_InitNewCycle_(PCPUACOSchedVars &pcpu_sched_vars) { pcpu_sched_vars.avlblSlotsInCrntCycle[i] = slotsPerTypePerCycle_[i]; } pcpu_sched_vars.isCrntCycleBlkd = false; +} + +void ACOScheduler::PCPU_InitSchedInsts(int numThreads){ + for (int i = 0; i < totInstCnt_; i++){ + SchedInstruction *inst = dataDepGraph_->GetInstByIndx(i); + inst->AllocPCPUVars(numThreads); + } +} + +void ACOScheduler::PCPU_FreeSchedInsts(int numThreads){ + for (int i = 0; i < totInstCnt_; i++){ + SchedInstruction *inst = dataDepGraph_->GetInstByIndx(i); + inst->FreePCPUVars(numThreads); + } } \ No newline at end of file diff --git a/lib/Scheduler/sched_basic_data.hip.cpp b/lib/Scheduler/sched_basic_data.hip.cpp index 8bef6bdc..87714965 100644 --- a/lib/Scheduler/sched_basic_data.hip.cpp +++ b/lib/Scheduler/sched_basic_data.hip.cpp @@ -1352,3 +1352,49 @@ __host__ __device__ bool SchedRange::IsTightnd(DIRECTION dir) const { return (dir == DIR_FRWRD) ? isFrwrdTightnd_ : isBkwrdTightnd_; } + +void SchedInstruction::AllocPCPUVars(int numThreads){ + pcpu_inst_vars_ = new PCPUSchedInstVars[numThreads]; + + for (int i = 0; i < numThreads; ++i) { + + pcpu_inst_vars_[i].crntSchedCycle = SCHD_UNSCHDULD; + pcpu_inst_vars_[i].crntSchedSlot = SCHD_UNSCHDULD; + pcpu_inst_vars_[i].lastUseCnt = 0; + pcpu_inst_vars_[i].ready = false; + pcpu_inst_vars_[i].minRdyCycle = INVALID_VALUE; + pcpu_inst_vars_[i].unschduldPrdcsrCnt = prdcsrCnt_; + pcpu_inst_vars_[i].unschduldScsrCnt = scsrCnt_; + pcpu_inst_vars_[i].crntRlxdCycle = SCHD_UNSCHDULD; + + if (prdcsrCnt_ > 0) { + pcpu_inst_vars_[i].rdyCyclePerPrdcsr = new InstCount[prdcsrCnt_]; + pcpu_inst_vars_[i].prevMinRdyCyclePerPrdcsr = new InstCount[prdcsrCnt_]; + + for (InstCount j = 0; j < prdcsrCnt_; ++j) { + pcpu_inst_vars_[i].rdyCyclePerPrdcsr[j] = INVALID_VALUE; + pcpu_inst_vars_[i].prevMinRdyCyclePerPrdcsr[j] = INVALID_VALUE; + } + } else { + pcpu_inst_vars_[i].rdyCyclePerPrdcsr = nullptr; + pcpu_inst_vars_[i].prevMinRdyCyclePerPrdcsr = nullptr; + } + } +} + +void SchedInstruction::FreePCPUVars(int numThreads) { + if (!pcpu_inst_vars_) + return; + + for (int i = 0; i < numThreads; ++i) { + delete[] pcpu_inst_vars_[i].rdyCyclePerPrdcsr; + delete[] pcpu_inst_vars_[i].prevMinRdyCyclePerPrdcsr; + + pcpu_inst_vars_[i].rdyCyclePerPrdcsr = nullptr; + pcpu_inst_vars_[i].prevMinRdyCyclePerPrdcsr = nullptr; + } + + delete[] pcpu_inst_vars_; + pcpu_inst_vars_ = nullptr; +} + From efd4c1355c1a492159e0167e1158b20156c731a1 Mon Sep 17 00:00:00 2001 From: Steven Inouye Date: Sat, 25 Apr 2026 22:33:53 -0700 Subject: [PATCH 17/30] Edited: Applied instruction functions, stil broken --- .../opt-sched/Scheduler/sched_basic_data.h | 8 +++ lib/Scheduler/aco.hip.cpp | 14 ++--- lib/Scheduler/sched_basic_data.hip.cpp | 51 +++++++++++++++++++ 3 files changed, 66 insertions(+), 7 deletions(-) diff --git a/include/opt-sched/Scheduler/sched_basic_data.h b/include/opt-sched/Scheduler/sched_basic_data.h index e88476da..e16177c9 100644 --- a/include/opt-sched/Scheduler/sched_basic_data.h +++ b/include/opt-sched/Scheduler/sched_basic_data.h @@ -581,6 +581,14 @@ class SchedInstruction : public GraphNode { void AllocPCPUVars(int numThreads); void InitPCPUVars(int numThreads); void FreePCPUVars(int numThreds); + bool PCPU_PrdcsrSchduld(InstCount prdcsrNum, InstCount cycle, + InstCount &rdyCycle, + int tIdx); + void PCPU_Schedule(InstCount cycleNum, InstCount slotNum, int tIdx); + int16_t PCPU_CmputLastUseCnt(RegisterFile *RegFiles, + int tIdx, + DataDepGraph *ddg = NULL); + int16_t PCPU_GetLastUseCnt(int tIdx); friend class SchedRange; diff --git a/lib/Scheduler/aco.hip.cpp b/lib/Scheduler/aco.hip.cpp index d6ca0b7d..83faf5d7 100644 --- a/lib/Scheduler/aco.hip.cpp +++ b/lib/Scheduler/aco.hip.cpp @@ -55,7 +55,7 @@ double RandDouble(double min, double max) { //#endif #define RUN_PCPU 1 -#define NO_CPU_THREADS 2 +#define NO_CPU_THREADS 1 ACOScheduler::ACOScheduler(DataDepGraph *dataDepGraph, MachineModel *machineModel, InstCount upperBound, @@ -2512,7 +2512,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, if (*pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) == pcpu_sched_vars.crntCycleNum) { InstCount CandidateId = *pcpu_sched_vars.readyLs->getInstIdAtIndex(I); SchedInstruction *candidateInst = dataDepGraph_->GetInstByIndx(CandidateId); - HeurType candidateLUC = candidateInst->GetLastUseCnt(); + HeurType candidateLUC = candidateInst->PCPU_GetLastUseCnt(thread); int16_t candidateDefs = candidateInst->GetDefCnt(); if (candidateDefs <= candidateLUC) { pcpu_sched_vars.RP0OrPositiveCount = pcpu_sched_vars.RP0OrPositiveCount + 1; @@ -2584,7 +2584,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, instNum = inst->GetNum(); PCPU_SchdulInst_(inst, pcpu_sched_vars); //SchdulInst(inst, pcpu_sched_vars.crntCycleNum) - inst->Schedule(pcpu_sched_vars.crntCycleNum, pcpu_sched_vars.crntSlotNum); + inst->PCPU_Schedule(pcpu_sched_vars.crntCycleNum, pcpu_sched_vars.crntSlotNum, thread); ((BBWithSpill *)rgn_)->PCPU_SchdulInst(inst, pcpu_sched_vars.crntCycleNum, pcpu_sched_vars.crntSlotNum, false, pcpu, thread); // If an ant violates the RP cost constraint, terminate further // schedule construction @@ -2645,7 +2645,7 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC RPIsHigh = false; InstCount CandidateId = *pcpu_sched_vars.readyLs->getInstIdAtIndex(I); SchedInstruction *candidateInst = dataDepGraph_->GetInstByIndx(CandidateId); - HeurType candidateLUC = candidateInst->GetLastUseCnt(); + HeurType candidateLUC = candidateInst->PCPU_GetLastUseCnt(thread); int16_t candidateDefs = candidateInst->GetDefCnt(); // compute the score @@ -2773,7 +2773,7 @@ inline void ACOScheduler::PCPU_UpdateACOReadyList(SchedInstruction *inst, bool I for (SchedInstruction *crntScsr = inst->GetFrstScsr(&prdcsrNum); crntScsr != NULL; crntScsr = inst->GetNxtScsr(&prdcsrNum)) { bool wasLastPrdcsr = - crntScsr->PrdcsrSchduld(prdcsrNum, pcpu_sched_vars.crntCycleNum, scsrRdyCycle); + crntScsr->PCPU_PrdcsrSchduld(prdcsrNum, pcpu_sched_vars.crntCycleNum, scsrRdyCycle, thread); if (wasLastPrdcsr) { // If all other predecessors of this successor have been scheduled then @@ -2794,7 +2794,7 @@ inline void ACOScheduler::PCPU_UpdateACOReadyList(SchedInstruction *inst, bool I InstCount CandidateId = *pcpu_sched_vars.readyLs->getInstIdAtIndex(I); if (LUCEntry.Width) { SchedInstruction *ScsrInst = dataDepGraph_->GetInstByIndx(CandidateId); - HeurType LUCVal = ScsrInst->CmputLastUseCnt(dataDepGraph_->RegFiles); + HeurType LUCVal = ScsrInst->PCPU_CmputLastUseCnt(dataDepGraph_->RegFiles, thread); LUCVal <<= LUCEntry.Offset; Heur &= LUCVal; } @@ -2803,7 +2803,7 @@ inline void ACOScheduler::PCPU_UpdateACOReadyList(SchedInstruction *inst, bool I continue; SchedInstruction *candidateInst = dataDepGraph_->GetInstByIndx(CandidateId); - HeurType candidateLUC = candidateInst->GetLastUseCnt(); + HeurType candidateLUC = candidateInst->PCPU_GetLastUseCnt(thread); int16_t candidateDefs = candidateInst->GetDefCnt(); if (candidateDefs <= candidateLUC) { pcpu_sched_vars.RP0OrPositiveCount = pcpu_sched_vars.RP0OrPositiveCount + 1; diff --git a/lib/Scheduler/sched_basic_data.hip.cpp b/lib/Scheduler/sched_basic_data.hip.cpp index 87714965..aeaa698d 100644 --- a/lib/Scheduler/sched_basic_data.hip.cpp +++ b/lib/Scheduler/sched_basic_data.hip.cpp @@ -1398,3 +1398,54 @@ void SchedInstruction::FreePCPUVars(int numThreads) { pcpu_inst_vars_ = nullptr; } +bool SchedInstruction::PCPU_PrdcsrSchduld(InstCount prdcsrNum, InstCount cycle, + InstCount &rdyCycle, + int tIdx ){ + assert(prdcsrNum < prdcsrCnt_); + assert(pcpu_inst_vars_[tIdx].rdyCyclePerPrdcsr != nullptr); + + pcpu_inst_vars_[tIdx].rdyCyclePerPrdcsr[prdcsrNum] = cycle + ltncyPerPrdcsr_[prdcsrNum]; + pcpu_inst_vars_[tIdx].prevMinRdyCyclePerPrdcsr[prdcsrNum] = pcpu_inst_vars_[tIdx].minRdyCycle; + + if (pcpu_inst_vars_[tIdx].rdyCyclePerPrdcsr[prdcsrNum] > pcpu_inst_vars_[tIdx].minRdyCycle) { + pcpu_inst_vars_[tIdx].minRdyCycle = pcpu_inst_vars_[tIdx].rdyCyclePerPrdcsr[prdcsrNum]; + } + + rdyCycle = pcpu_inst_vars_[tIdx].minRdyCycle; + pcpu_inst_vars_[tIdx].unschduldPrdcsrCnt--; + return (pcpu_inst_vars_[tIdx].unschduldPrdcsrCnt == 0); +} + +void SchedInstruction::PCPU_Schedule(InstCount cycleNum, InstCount slotNum, + int tIdx){ + assert(pcpu_inst_vars_[tIdx].crntSchedCycle == SCHD_UNSCHDULD); + pcpu_inst_vars_[tIdx].crntSchedCycle = cycleNum; + pcpu_inst_vars_[tIdx].crntSchedSlot = slotNum; +} + +int16_t SchedInstruction::PCPU_CmputLastUseCnt(RegisterFile *RegFiles, + int tIdx, + DataDepGraph *ddg) { + pcpu_inst_vars_[tIdx].lastUseCnt = 0; + + // if we are on device or passing in RegFiles, use those + // otherwise, assume we can get regFiles from SchedInstruction itself + RegisterFile *registerFiles; + if (RegFiles) + registerFiles = RegFiles; + else + registerFiles = RegFiles_; + + for (int i = 0; i < useCnt_; i++) { + Register *reg = registerFiles[uses_[i].regType_].GetReg(uses_[i].regNum_); + assert(reg->GetCrntUseCnt() < reg->GetUseCnt()); + if (reg->GetCrntUseCnt() + 1 == reg->GetUseCnt()) + pcpu_inst_vars_[tIdx].lastUseCnt++; + } + return pcpu_inst_vars_[tIdx].lastUseCnt; +} + +int16_t SchedInstruction::PCPU_GetLastUseCnt(int tIdx) { + return pcpu_inst_vars_[tIdx].lastUseCnt; +} + From 138fb2d43d70eff90b7346e30608884c70ef5353 Mon Sep 17 00:00:00 2001 From: Steven Inouye Date: Sun, 26 Apr 2026 10:44:24 -0700 Subject: [PATCH 18/30] temp --- include/opt-sched/Scheduler/aco.h | 6 +----- lib/Scheduler/aco.hip.cpp | 9 +++++---- 2 files changed, 6 insertions(+), 9 deletions(-) diff --git a/include/opt-sched/Scheduler/aco.h b/include/opt-sched/Scheduler/aco.h index e8da2e51..cd99d63a 100644 --- a/include/opt-sched/Scheduler/aco.h +++ b/include/opt-sched/Scheduler/aco.h @@ -115,12 +115,7 @@ struct alignas(64) PCPUACOSchedVars { isCrntCycleBlkd_ avlblSlotsInCrntCycle_ */ - - - - //may need multiple khelpers as well unsure - }; class ACOScheduler : public ConstrainedScheduler { @@ -163,6 +158,7 @@ class ACOScheduler : public ConstrainedScheduler { bool closeToRPTarget, bool currentlyWaiting, int thread, PCPUACOSchedVars &pcpu_sched_vars, + ParallelCPUVars &pcpu, int kernelNum = -1); inline void PCPU_UpdateACOReadyList(SchedInstruction *inst, bool IsSecondPass, int thread, diff --git a/lib/Scheduler/aco.hip.cpp b/lib/Scheduler/aco.hip.cpp index 83faf5d7..b08c8992 100644 --- a/lib/Scheduler/aco.hip.cpp +++ b/lib/Scheduler/aco.hip.cpp @@ -2448,7 +2448,7 @@ InstSchedule *ACOScheduler::FindManyCPUSchedule(InstCount RPTarget) { } ((BBWithSpill*)rgn_)->FreeParallelCPUVars(NO_CPU_THREADS); - //logic to find best schedule + //logic to find best schedule should be added here InstSchedule *result = cpuScheds[0]; for (int i = 1; i < NO_CPU_THREADS; i++) { @@ -2622,6 +2622,7 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC bool closeToRPTarget, bool currentlyWaiting, int thread, PCPUACOSchedVars &pcpu_sched_vars, + ParallelCPUVars &pcpu, int blockOccupancyNum){ // if we are waiting and have no fully-ready instruction that is // net 0 or benefit to RP, then return -1 to schedule a stall @@ -2688,7 +2689,7 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC for (uint16_t i = 0; i < usesCount; i++) { use = dataDepGraph_->getRegByTuple(&uses[i]); int16_t regType = use->GetType(); - if ( ((BBWithSpill *)rgn)->IsRPHigh(regType) ) { + if ( ((BBWithSpill *)rgn)->PCPU_IsRPHigh(regType, pcpu) ) { RPIsHigh = true; break; } @@ -2870,9 +2871,9 @@ void ACOScheduler::PCPU_DoRsrvSlots_(SchedInstruction *inst, PCPUACOSchedVars &p return; if (!inst->IsPipelined()) { - if (pcpu_sched_vars.rsrvSlots == NULL){ + /* if (pcpu_sched_vars.rsrvSlots == NULL){ //AllocRsrvSlots_(); //not sure if this function is necessary - } + } */ pcpu_sched_vars.rsrvSlots[pcpu_sched_vars.crntSlotNum].strtCycle = pcpu_sched_vars.crntCycleNum; pcpu_sched_vars.rsrvSlots[pcpu_sched_vars.crntSlotNum].endCycle = pcpu_sched_vars.crntCycleNum + inst->GetMaxLtncy() - 1; pcpu_sched_vars.rsrvSlotCnt++; From 48f5e08b8f3b5fb150f09a5cd136d742d7984c8a Mon Sep 17 00:00:00 2001 From: Steven Inouye Date: Sun, 3 May 2026 13:58:21 -0700 Subject: [PATCH 19/30] temp --- lib/Scheduler/aco.hip.cpp | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) diff --git a/lib/Scheduler/aco.hip.cpp b/lib/Scheduler/aco.hip.cpp index b08c8992..0e1bca1e 100644 --- a/lib/Scheduler/aco.hip.cpp +++ b/lib/Scheduler/aco.hip.cpp @@ -2005,7 +2005,7 @@ void ACOScheduler::UpdatePheromone(InstSchedule *schedule, bool isIterationBest, #if !USE_ACS // decay pheromone for (int i = 0; i < count_; i++) { - for (int j = 0; j < count_; j++) { + for (int j = 0; j < count_; j++) { pheromone = &Pheromone(i, j); *pheromone *= (1 - decay_factor); } @@ -2532,7 +2532,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, closeToRPTarget = ((BBWithSpill *)rgn_)->PCPU_GetCrntSpillCost(pcpu) >= closeToRPCheck; // select the instruction and get info on it InstCount SelIndx = PCPU_SelectInstruction(lastInst, schedule->getTotalStalls(), rgn_, unnecessarilyStalling, closeToRPTarget, waitFor ? true: false, - thread, pcpu_sched_vars); + thread, pcpu_sched_vars, pcpu); if (SelIndx != -1) { LastInstInfo = pcpu_sched_vars.readyLs->removeInstructionAtIndex(SelIndx); @@ -2761,6 +2761,7 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC else unnecessarilyStalling = false; + Logger::Info("Instruction Index, %d", indx); return indx; } @@ -2912,6 +2913,7 @@ void ACOScheduler::PCPU_SchdulInst_(SchedInstruction *inst, PCPUACOSchedVars &pc pcpu_sched_vars.isCrntCycleBlkd = true; } pcpu_sched_vars.schduldInstCnt++; + Logger::Info("SchduldInstCnt Incremented, %d", pcpu_sched_vars.schduldInstCnt); } void ACOScheduler::PCPU_UpdtSlotAvlblty_(SchedInstruction *inst, PCPUACOSchedVars &pcpu_sched_vars) { @@ -2925,6 +2927,7 @@ void ACOScheduler::PCPU_UpdtSlotAvlblty_(SchedInstruction *inst, PCPUACOSchedVar } bool ACOScheduler::PCPU_IsSchedComplete_(PCPUACOSchedVars &pcpu_sched_vars) { + Logger::Info("Checking if sched complete if %d == %d", pcpu_sched_vars.schduldInstCnt, totInstCnt_); return pcpu_sched_vars.schduldInstCnt == totInstCnt_; } From 498ea4d00160104eef71a5d9f284259e0d8b88e9 Mon Sep 17 00:00:00 2001 From: Steven Inouye Date: Sat, 9 May 2026 14:33:44 -0700 Subject: [PATCH 20/30] Added: Accurate MaxPriorityInv and commeted out number of ants terminated adjustment. Problem remains. --- include/opt-sched/Scheduler/aco.h | 1 + lib/Scheduler/aco.hip.cpp | 37 +++++++++++++++++++------------ 2 files changed, 24 insertions(+), 14 deletions(-) diff --git a/include/opt-sched/Scheduler/aco.h b/include/opt-sched/Scheduler/aco.h index cd99d63a..8403b1e8 100644 --- a/include/opt-sched/Scheduler/aco.h +++ b/include/opt-sched/Scheduler/aco.h @@ -87,6 +87,7 @@ struct alignas(64) PCPUACOSchedVars { ReserveSlot *rsrvSlots; int16_t *avlblSlotsInCrntCycle; + pheromone_t MaxPriorityInv; /* void ConstrainedScheduler::SchdulInst_(SchedInstruction *inst, InstCount) diff --git a/lib/Scheduler/aco.hip.cpp b/lib/Scheduler/aco.hip.cpp index 0e1bca1e..a2a6526f 100644 --- a/lib/Scheduler/aco.hip.cpp +++ b/lib/Scheduler/aco.hip.cpp @@ -55,7 +55,7 @@ double RandDouble(double min, double max) { //#endif #define RUN_PCPU 1 -#define NO_CPU_THREADS 1 +#define NO_CPU_THREADS 2 ACOScheduler::ACOScheduler(DataDepGraph *dataDepGraph, MachineModel *machineModel, InstCount upperBound, @@ -1795,7 +1795,7 @@ FUNC_RESULT ACOScheduler::FindSchedule(InstSchedule *schedule_out, int diffSchedCount = 0; #endif - while (noImprovement < noImprovementMax) { + while (noImprovement < noImprovementMax) { iterations++; iterationBest = nullptr; for (int i = 0; i < numThreads_; i++) { @@ -1804,6 +1804,7 @@ FUNC_RESULT ACOScheduler::FindSchedule(InstSchedule *schedule_out, schedule = FindOneSchedule(RPTarget, NULL); } else { schedule = FindManyCPUSchedule(RPTarget); + Logger::Info("Done Find Many Schedule"); i+=(NO_CPU_THREADS-1); } @@ -2450,6 +2451,9 @@ InstSchedule *ACOScheduler::FindManyCPUSchedule(InstCount RPTarget) { //logic to find best schedule should be added here InstSchedule *result = cpuScheds[0]; + MaxPriorityInv = pcpu_sched_vars[0].MaxPriorityInv; + //the index with the best schedule should also grab the pcpu_sched_vars[i].MaxPriorityInv + //and update the global MaxPriorityInv for (int i = 1; i < NO_CPU_THREADS; i++) { delete cpuScheds[i]; @@ -2460,7 +2464,7 @@ InstSchedule *ACOScheduler::FindManyCPUSchedule(InstCount RPTarget) { delete[] cpuScheds; PCPU_FreeSchedInsts(NO_CPU_THREADS); - //printf("Ending FindManyCPUSchedule"); + printf("Ending FindManyCPUSchedule"); return result; } @@ -2485,7 +2489,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, SchedInstruction *waitFor = NULL; InstCount waitUntil = 0; - MaxPriorityInv = 1 / (pheromone_t)MaxPriority; + pcpu_sched_vars.MaxPriorityInv = 1 / (pheromone_t)MaxPriority; // initialize the aco ready list so that the start instruction is ready // The luc component is 0 since the root inst uses no instructions @@ -2507,7 +2511,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, // incrementally calculate if there are any instructions with a neutral // or positive effect on RP - Logger::Info("Neut/pos on RP, %d", thread); + //Logger::Info("Neut/pos on RP, %d", thread); for (InstCount I = 0; I < pcpu_sched_vars.readyLs->getReadyListSize(); ++I) { if (*pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) == pcpu_sched_vars.crntCycleNum) { InstCount CandidateId = *pcpu_sched_vars.readyLs->getInstIdAtIndex(I); @@ -2522,7 +2526,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, // there are two steps to scheduling an instruction: // 1)Select the instruction(if we are not waiting on another instruction) - Logger::Info("SelectInstruction, %d", thread); + //Logger::Info("SelectInstruction, %d", thread); inst = NULL; if (!(waitFor && waitUntil <= pcpu_sched_vars.crntCycleNum)) { // If an instruction is ready select it @@ -2561,7 +2565,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, } } - Logger::Info("Scheduling Stalls, %d", thread); + //Logger::Info("Scheduling Stalls, %d", thread); // 2)Schedule a stall if we are still waiting, Schedule the instruction we // are waiting for if possible, decrement waiting time if (waitFor && waitUntil <= pcpu_sched_vars.crntCycleNum) { @@ -2573,7 +2577,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, } // boilerplate, mostly copied from ListScheduler, try not to touch it - Logger::Info("BoilerPlate, %d", thread); + //Logger::Info("BoilerPlate, %d", thread); InstCount instNum; if (!inst) { instNum = SCHD_STALL; @@ -2594,6 +2598,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, numAntsTerminated_++; pcpu_sched_vars.readyLs->clearReadyList(); delete schedule; + Logger::Info("Thread %d, returned early", thread); return NULL; } PCPU_DoRsrvSlots_(inst, pcpu_sched_vars); @@ -2613,7 +2618,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, ((BBWithSpill *)rgn_)->PCPU_UpdateScheduleCost(schedule, pcpu); schedule->setIsZeroPerp(((BBWithSpill *)rgn_)->PCPU_ReturnPeakSpillCost(pcpu) == 0 ); schedule->setOccupancy(((BBWithSpill *)rgn_)->PCPU_getOccupancy(pcpu)); - Logger::Info("Done, %d", thread); + Logger::Info("Thread, %d, Done", thread); return schedule; } @@ -2761,16 +2766,17 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC else unnecessarilyStalling = false; - Logger::Info("Instruction Index, %d", indx); + //Logger::Info("Instruction Index, %d", indx); return indx; } +//culprit im thinking could be the issue inline void ACOScheduler::PCPU_UpdateACOReadyList(SchedInstruction *inst, bool IsSecondPass, int thread, PCPUACOSchedVars &pcpu_sched_vars, int heurChoice){ InstCount prdcsrNum, scsrRdyCycle; - Logger::Info("Start UpdateACOReadyList, %d", thread); + //Logger::Info("Start UpdateACOReadyList, %d", thread); // Notify each successor of this instruction that it has been scheduled. for (SchedInstruction *crntScsr = inst->GetFrstScsr(&prdcsrNum); crntScsr != NULL; crntScsr = inst->GetNxtScsr(&prdcsrNum)) { @@ -2812,7 +2818,7 @@ inline void ACOScheduler::PCPU_UpdateACOReadyList(SchedInstruction *inst, bool I } } } - Logger::Info("Finish UpdateACOReadyList, %d", thread); + //Logger::Info("Finish UpdateACOReadyList, %d", thread); } //allocates memory and intializes variables @@ -2831,6 +2837,9 @@ PCPUACOSchedVars *ACOScheduler::AllocPCPUACOSchedVars(int numThreads) { pcpu_sched_vars[i].crntSlotNum = 0; pcpu_sched_vars[i].rsrvSlotCnt = 0; + //new stuff + pcpu_sched_vars[i].MaxPriorityInv = 0; + pcpu_sched_vars[i].rsrvSlots = new ReserveSlot[issuRate_]; pcpu_sched_vars[i].avlblSlotsInCrntCycle = new int16_t[issuTypeCnt_]; @@ -2913,7 +2922,7 @@ void ACOScheduler::PCPU_SchdulInst_(SchedInstruction *inst, PCPUACOSchedVars &pc pcpu_sched_vars.isCrntCycleBlkd = true; } pcpu_sched_vars.schduldInstCnt++; - Logger::Info("SchduldInstCnt Incremented, %d", pcpu_sched_vars.schduldInstCnt); + //Logger::Info("SchduldInstCnt Incremented, %d", pcpu_sched_vars.schduldInstCnt); } void ACOScheduler::PCPU_UpdtSlotAvlblty_(SchedInstruction *inst, PCPUACOSchedVars &pcpu_sched_vars) { @@ -2927,7 +2936,7 @@ void ACOScheduler::PCPU_UpdtSlotAvlblty_(SchedInstruction *inst, PCPUACOSchedVar } bool ACOScheduler::PCPU_IsSchedComplete_(PCPUACOSchedVars &pcpu_sched_vars) { - Logger::Info("Checking if sched complete if %d == %d", pcpu_sched_vars.schduldInstCnt, totInstCnt_); + //Logger::Info("Checking if sched complete if %d == %d", pcpu_sched_vars.schduldInstCnt, totInstCnt_); return pcpu_sched_vars.schduldInstCnt == totInstCnt_; } From 9591486a64b6207eafc21ebbec3e554cecd507e7 Mon Sep 17 00:00:00 2001 From: Steven Inouye Date: Sat, 9 May 2026 14:33:54 -0700 Subject: [PATCH 21/30] temp; --- lib/Scheduler/aco.hip.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/lib/Scheduler/aco.hip.cpp b/lib/Scheduler/aco.hip.cpp index a2a6526f..1a53b17c 100644 --- a/lib/Scheduler/aco.hip.cpp +++ b/lib/Scheduler/aco.hip.cpp @@ -2595,7 +2595,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, if (((BBWithSpill*)rgn_)->PCPU_GetCrntSpillCost(pcpu) > RPTarget) { // end schedule construction // keep track of ants terminated - numAntsTerminated_++; + //numAntsTerminated_++; pcpu_sched_vars.readyLs->clearReadyList(); delete schedule; Logger::Info("Thread %d, returned early", thread); From c9728a457f4665cee1762e9c8a3fcafecfcbf99c Mon Sep 17 00:00:00 2001 From: Steven Inouye Date: Sun, 10 May 2026 10:42:34 -0700 Subject: [PATCH 22/30] Added: New score function, errors remain. --- include/opt-sched/Scheduler/aco.h | 4 ++++ lib/Scheduler/aco.hip.cpp | 29 ++++++++++++++++++++++------- 2 files changed, 26 insertions(+), 7 deletions(-) diff --git a/include/opt-sched/Scheduler/aco.h b/include/opt-sched/Scheduler/aco.h index 8403b1e8..b6ed43fb 100644 --- a/include/opt-sched/Scheduler/aco.h +++ b/include/opt-sched/Scheduler/aco.h @@ -88,6 +88,7 @@ struct alignas(64) PCPUACOSchedVars { int16_t *avlblSlotsInCrntCycle; pheromone_t MaxPriorityInv; + KeysHelper1 *kHelper1; /* void ConstrainedScheduler::SchdulInst_(SchedInstruction *inst, InstCount) @@ -230,6 +231,9 @@ class ACOScheduler : public ConstrainedScheduler { pheromone_t &Pheromone(InstCount from, InstCount to, int kernelNum = 0); __host__ __device__ pheromone_t Score(InstCount FromId, InstCount ToId, HeurType ToHeuristic, bool IsFirstPass, int kernelNum = 0); + + pheromone_t PCPU_Score(InstCount FromId, InstCount ToId, HeurType ToHeuristic, bool IsFirstPass, PCPUACOSchedVars &pcpu_sched_vars, int kernelNum = 0); + DCF_OPT ParseDCFOpt(const std::string &opt); __host__ __device__ InstCount SelectInstruction(SchedInstruction *lastInst, InstCount totalStalls, diff --git a/lib/Scheduler/aco.hip.cpp b/lib/Scheduler/aco.hip.cpp index 1a53b17c..cc3490fe 100644 --- a/lib/Scheduler/aco.hip.cpp +++ b/lib/Scheduler/aco.hip.cpp @@ -55,7 +55,7 @@ double RandDouble(double min, double max) { //#endif #define RUN_PCPU 1 -#define NO_CPU_THREADS 2 +#define NO_CPU_THREADS 1 ACOScheduler::ACOScheduler(DataDepGraph *dataDepGraph, MachineModel *machineModel, InstCount upperBound, @@ -2482,7 +2482,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, bool unnecessarilyStalling = false; // The MaxPriority that we are getting from the ready list represents the maximum possible heuristic/key value that we can have // I want to move all the heuristic computation stuff to another class for code tidiness reasons. - HeurType MaxPriority = kHelper1->getMaxValue(); + HeurType MaxPriority = pcpu_sched_vars.kHelper1->getMaxValue(); if (MaxPriority == 0) MaxPriority = 1; // divide by 0 is bad //Initialize_(); //For InstSchedule/ConstrainedSchedule @@ -2495,8 +2495,8 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, // The luc component is 0 since the root inst uses no instructions Logger::Info("ReadyListFunctions 1, %d", thread); InstCount RootId = rootInst_->GetNum(); - HeurType RootHeuristic = kHelper1->computeKey(rootInst_, true, dataDepGraph_->RegFiles); - pheromone_t RootScore = Score(-1, RootId, RootHeuristic, !IsSecondPass); + HeurType RootHeuristic = pcpu_sched_vars.kHelper1->computeKey(rootInst_, true, dataDepGraph_->RegFiles); + pheromone_t RootScore = PCPU_Score(-1, RootId, RootHeuristic, !IsSecondPass, pcpu_sched_vars); ACOReadyListEntry InitialRoot{RootId, 0, RootHeuristic, RootScore}; pcpu_sched_vars.readyLs->addInstructionToReadyList(InitialRoot); pcpu_sched_vars.readyLs->ScoreSum = RootScore; @@ -2656,7 +2656,7 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC // compute the score HeurType Heur = *pcpu_sched_vars.readyLs->getInstHeuristicAtIndex(I); - pheromone_t IScore = Score(lastInstId, *pcpu_sched_vars.readyLs->getInstIdAtIndex(I), Heur, !rgn->IsSecondPass()); + pheromone_t IScore = PCPU_Score(lastInstId, *pcpu_sched_vars.readyLs->getInstIdAtIndex(I), Heur, !rgn->IsSecondPass(), pcpu_sched_vars); if (pcpu_sched_vars.RP0OrPositiveCount != 0 && candidateDefs > candidateLUC) IScore = IScore * 9/10; @@ -2786,14 +2786,14 @@ inline void ACOScheduler::PCPU_UpdateACOReadyList(SchedInstruction *inst, bool I if (wasLastPrdcsr) { // If all other predecessors of this successor have been scheduled then // we now know in which cycle this successor will become ready. - HeurType HeurWOLuc = kHelper1->computeKey(crntScsr, false, dataDepGraph_->RegFiles); + HeurType HeurWOLuc = pcpu_sched_vars.kHelper1->computeKey(crntScsr, false, dataDepGraph_->RegFiles); pcpu_sched_vars.readyLs->addInstructionToReadyList(ACOReadyListEntry{crntScsr->GetNum(), scsrRdyCycle, HeurWOLuc, 0}); } } // Make sure the scores are valid. The scheduling of an instruction may // have increased another instruction's LUC Score - PriorityEntry LUCEntry = kHelper1->getPriorityEntry(LSH_LUC); + PriorityEntry LUCEntry = pcpu_sched_vars.kHelper1->getPriorityEntry(LSH_LUC); pcpu_sched_vars.RP0OrPositiveCount = 0; for (InstCount I = 0; I < pcpu_sched_vars.readyLs->getReadyListSize(); ++I) { //we first get the heuristic without the LUC component, add the LUC @@ -2839,6 +2839,8 @@ PCPUACOSchedVars *ACOScheduler::AllocPCPUACOSchedVars(int numThreads) { //new stuff pcpu_sched_vars[i].MaxPriorityInv = 0; + pcpu_sched_vars[i].kHelper1 = new KeysHelper1(priorities1_); + pcpu_sched_vars[i].kHelper1->initForRegion(dataDepGraph_); pcpu_sched_vars[i].rsrvSlots = new ReserveSlot[issuRate_]; pcpu_sched_vars[i].avlblSlotsInCrntCycle = new int16_t[issuTypeCnt_]; @@ -2865,6 +2867,9 @@ void ACOScheduler::FreePCPUACOSchedVars(PCPUACOSchedVars *pcpu_sched_vars, delete pcpu_sched_vars[i].readyLs; pcpu_sched_vars[i].readyLs = nullptr; + delete pcpu_sched_vars[i].kHelper1; + pcpu_sched_vars[i].kHelper1 = nullptr; + delete[] pcpu_sched_vars[i].rsrvSlots; pcpu_sched_vars[i].rsrvSlots = nullptr; @@ -3004,4 +3009,14 @@ void ACOScheduler::PCPU_FreeSchedInsts(int numThreads){ SchedInstruction *inst = dataDepGraph_->GetInstByIndx(i); inst->FreePCPUVars(numThreads); } +} + + +pheromone_t ACOScheduler::PCPU_Score(InstCount FromId, InstCount ToId, HeurType ToHeuristic, bool IsFirstPass, PCPUACOSchedVars &pcpu_sched_vars, int blockOccupancyNum) { + // tuneable heuristic importance is temporarily disabled + // double Hf = pow(ToHeuristic, heuristicImportance_); + pheromone_t HeurScore; + HeurScore = ToHeuristic * pcpu_sched_vars.MaxPriorityInv + 1; + pheromone_t Hf = heuristicImportance_ ? HeurScore : 1.0; + return Pheromone(FromId, ToId, blockOccupancyNum) * Hf; } \ No newline at end of file From f34c9d84294631ebd0022677b301904427b1c085 Mon Sep 17 00:00:00 2001 From: mocha8686 Date: Sat, 9 May 2026 16:53:34 -0700 Subject: [PATCH 23/30] fix: add recursive_mutex to random number generator --- lib/Scheduler/random.hip.cpp | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/lib/Scheduler/random.hip.cpp b/lib/Scheduler/random.hip.cpp index c998b460..93d34c9c 100644 --- a/lib/Scheduler/random.hip.cpp +++ b/lib/Scheduler/random.hip.cpp @@ -2,9 +2,13 @@ // For memcpy(). #include #include +#include +#include using namespace llvm::opt_sched; +static const std::recursive_mutex m; + // Magic numbers used in the generator formula. static const uint32_t A = 0x2faf071d; // 8 * (10 ** 8 - 29) + 5 static const uint32_t C = 0x3b9ac9c1; // 10 ** 9 - 63 @@ -43,6 +47,7 @@ static uint32_t y[] = { static uint32_t randNum; void GenerateNextNumber() { + std::scoped_lock{m}; randNum = y[j] + y[k]; y[k] = randNum; if (--j < 0) @@ -53,6 +58,7 @@ void GenerateNextNumber() { } void RandomGen::SetSeed(int32_t iseed) { + std::scoped_lock{m}; j = 23; k = 54; @@ -69,16 +75,19 @@ void RandomGen::SetSeed(int32_t iseed) { } uint32_t RandomGen::GetRand32WithinRange(uint32_t min, uint32_t max) { + std::scoped_lock{m}; GenerateNextNumber(); return randNum % (max - min + 1) + min; } uint32_t RandomGen::GetRand32() { + std::scoped_lock{m}; GenerateNextNumber(); return randNum; } uint64_t RandomGen::GetRand64() { + std::scoped_lock{m}; uint64_t rand64; GenerateNextNumber(); @@ -92,6 +101,7 @@ uint64_t RandomGen::GetRand64() { } void RandomGen::GetRandBits(uint16_t bitCnt, unsigned char *dest) { + std::scoped_lock{m}; uint16_t bytesNeeded = (bitCnt + 7) / 8; uint16_t index = 0; From 30bd74ff05c943ef73870ee4ccafe9909ff03df2 Mon Sep 17 00:00:00 2001 From: mocha8686 Date: Sat, 9 May 2026 16:56:06 -0700 Subject: [PATCH 24/30] fix: remove uninitialized int usages and add default edge labels to 0 --- include/opt-sched/Scheduler/graph.h | 28 ++++------------------------ 1 file changed, 4 insertions(+), 24 deletions(-) diff --git a/include/opt-sched/Scheduler/graph.h b/include/opt-sched/Scheduler/graph.h index a1539020..9249d735 100644 --- a/include/opt-sched/Scheduler/graph.h +++ b/include/opt-sched/Scheduler/graph.h @@ -159,23 +159,15 @@ class GraphNode { __host__ void AddRcrsvNghbr(GraphNode *nghbr, DIRECTION dir); // Returns a pointer to the first successor of the node and writes the label - // of the edge between them to the label argument. Sets the successor + // of the edge between them to the label argument (default 0). Sets the successor // iterator. __host__ GraphNode *GetFrstScsr(UDT_GLABEL &label); // Returns a pointer to the next successor of the node and writes the label - // of the edge between them to the label argument. Must be called after + // of the edge between them to the label argument (default 0). Must be called after // GetFrstScsr() which starts the successor iterator. __host__ GraphNode *GetNxtScsr(UDT_GLABEL &label); - // Returns a pointer to the first successor of the node. Sets the successor - // iterator. - __host__ - GraphNode *GetFrstScsr(); - // Returns a pointer to the next successor of the node. Must be called after - // GetFrstScsr() which starts the successor iterator. - __host__ - GraphNode *GetNxtScsr(); // Checks if a given node is successor-equivalent to this node. Two nodes // are successor-equivalent if they have identical successor lists. __host__ @@ -591,7 +583,7 @@ __host__ __device__ inline UDT_GNODES GraphNode::GetNum() const { return num_; } __host__ -inline GraphNode *GraphNode::GetFrstScsr(UDT_GLABEL &label) { +inline GraphNode *GraphNode::GetFrstScsr(UDT_GLABEL &label = 0) { GraphEdge *edge = scsrLst_->GetFrstElmnt(); if (edge == NULL) return NULL; @@ -600,7 +592,7 @@ inline GraphNode *GraphNode::GetFrstScsr(UDT_GLABEL &label) { } __host__ -inline GraphNode *GraphNode::GetNxtScsr(UDT_GLABEL &label) { +inline GraphNode *GraphNode::GetNxtScsr(UDT_GLABEL &label = 0) { GraphEdge *edge = scsrLst_->GetNxtElmnt(); if (edge == NULL) return NULL; @@ -626,18 +618,6 @@ inline GraphNode *GraphNode::GetNxtPrdcsr(UDT_GLABEL &label) { return nodes_[edge->to]; } -__host__ -inline GraphNode *GraphNode::GetFrstScsr() { - UDT_GLABEL label; - return GetFrstScsr(label); -} - -__host__ -inline GraphNode *GraphNode::GetNxtScsr() { - UDT_GLABEL label; - return GetNxtScsr(label); -} - __host__ __device__ inline UDT_GLABEL GraphNode::GetScsrLblSum() const { return scsrLblSum_; } From 428e80584ce6ac89a8a2bf4c929c44c6c3cc000d Mon Sep 17 00:00:00 2001 From: mocha8686 Date: Sun, 10 May 2026 11:59:57 -0700 Subject: [PATCH 25/30] fix: cv-unqualify mutex --- lib/Scheduler/random.hip.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/lib/Scheduler/random.hip.cpp b/lib/Scheduler/random.hip.cpp index 93d34c9c..4c6d7228 100644 --- a/lib/Scheduler/random.hip.cpp +++ b/lib/Scheduler/random.hip.cpp @@ -7,7 +7,7 @@ using namespace llvm::opt_sched; -static const std::recursive_mutex m; +static std::recursive_mutex m; // Magic numbers used in the generator formula. static const uint32_t A = 0x2faf071d; // 8 * (10 ** 8 - 29) + 5 From 15a7cc33b51492d70258d509077c80da37744fc5 Mon Sep 17 00:00:00 2001 From: mocha8686 Date: Sun, 10 May 2026 12:06:04 -0700 Subject: [PATCH 26/30] fix: pass reference to function as default instead --- include/opt-sched/Scheduler/graph.h | 28 ++++++++++++++++++++++++---- 1 file changed, 24 insertions(+), 4 deletions(-) diff --git a/include/opt-sched/Scheduler/graph.h b/include/opt-sched/Scheduler/graph.h index 9249d735..8ca11726 100644 --- a/include/opt-sched/Scheduler/graph.h +++ b/include/opt-sched/Scheduler/graph.h @@ -159,15 +159,23 @@ class GraphNode { __host__ void AddRcrsvNghbr(GraphNode *nghbr, DIRECTION dir); // Returns a pointer to the first successor of the node and writes the label - // of the edge between them to the label argument (default 0). Sets the successor + // of the edge between them to the label argument. Sets the successor // iterator. __host__ GraphNode *GetFrstScsr(UDT_GLABEL &label); // Returns a pointer to the next successor of the node and writes the label - // of the edge between them to the label argument (default 0). Must be called after + // of the edge between them to the label argument. Must be called after // GetFrstScsr() which starts the successor iterator. __host__ GraphNode *GetNxtScsr(UDT_GLABEL &label); + // Returns a pointer to the first successor of the node. Sets the successor + // iterator. + __host__ + GraphNode *GetFrstScsr(); + // Returns a pointer to the next successor of the node. Must be called after + // GetFrstScsr() which starts the successor iterator. + __host__ + GraphNode *GetNxtScsr(); // Checks if a given node is successor-equivalent to this node. Two nodes // are successor-equivalent if they have identical successor lists. __host__ @@ -583,7 +591,7 @@ __host__ __device__ inline UDT_GNODES GraphNode::GetNum() const { return num_; } __host__ -inline GraphNode *GraphNode::GetFrstScsr(UDT_GLABEL &label = 0) { +inline GraphNode *GraphNode::GetFrstScsr(UDT_GLABEL &label) { GraphEdge *edge = scsrLst_->GetFrstElmnt(); if (edge == NULL) return NULL; @@ -592,7 +600,7 @@ inline GraphNode *GraphNode::GetFrstScsr(UDT_GLABEL &label = 0) { } __host__ -inline GraphNode *GraphNode::GetNxtScsr(UDT_GLABEL &label = 0) { +inline GraphNode *GraphNode::GetNxtScsr(UDT_GLABEL &label) { GraphEdge *edge = scsrLst_->GetNxtElmnt(); if (edge == NULL) return NULL; @@ -618,6 +626,18 @@ inline GraphNode *GraphNode::GetNxtPrdcsr(UDT_GLABEL &label) { return nodes_[edge->to]; } +__host__ +inline GraphNode *GraphNode::GetFrstScsr() { + UDT_GLABEL label{}; + return GetFrstScsr(label); +} + +__host__ +inline GraphNode *GraphNode::GetNxtScsr() { + UDT_GLABEL label{}; + return GetNxtScsr(label); +} + __host__ __device__ inline UDT_GLABEL GraphNode::GetScsrLblSum() const { return scsrLblSum_; } From 8366c5f779f2a99a3293f984a51fe48bdc1ab9c0 Mon Sep 17 00:00:00 2001 From: mocha8686 Date: Sun, 10 May 2026 19:28:37 -0700 Subject: [PATCH 27/30] fix: use 2 threads --- lib/Scheduler/aco.hip.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/lib/Scheduler/aco.hip.cpp b/lib/Scheduler/aco.hip.cpp index cc3490fe..8e9330ac 100644 --- a/lib/Scheduler/aco.hip.cpp +++ b/lib/Scheduler/aco.hip.cpp @@ -55,7 +55,7 @@ double RandDouble(double min, double max) { //#endif #define RUN_PCPU 1 -#define NO_CPU_THREADS 1 +#define NO_CPU_THREADS 2 ACOScheduler::ACOScheduler(DataDepGraph *dataDepGraph, MachineModel *machineModel, InstCount upperBound, @@ -3019,4 +3019,4 @@ pheromone_t ACOScheduler::PCPU_Score(InstCount FromId, InstCount ToId, HeurType HeurScore = ToHeuristic * pcpu_sched_vars.MaxPriorityInv + 1; pheromone_t Hf = heuristicImportance_ ? HeurScore : 1.0; return Pheromone(FromId, ToId, blockOccupancyNum) * Hf; -} \ No newline at end of file +} From 8b557bb178bc547138d11a8c9934ae34a0ae47f3 Mon Sep 17 00:00:00 2001 From: Peng Kang Date: Sat, 23 May 2026 00:35:56 +0000 Subject: [PATCH 28/30] fix: stabilize CPU-parallel ACO and switch to 8-thread execution MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit - random.hip.cpp: replace std::scoped_lock{m} (no-op temporary) with std::scoped_lock lock(m) in all six call sites; the anonymous temporary was immediately destroyed, leaving j/k/y[] unprotected and causing sporadic out-of-bounds SIGSEGV under concurrent RNG calls - aco.hip.cpp: pre-build instScsrs_ before spawning PCPU threads to eliminate the GetFrstScsr/GetNxtScsr cursor race on scsrLst_->rtrvEntry_ - aco.hip.cpp: fix integer division IScore*9/10→IScore*=0.9 and guard globalBestStalls_==0 before divisions in SelectInstruction and PCPU_SelectInstruction - aco.hip.cpp: select best schedule across all CPU threads in FindManyCPUSchedule instead of always taking thread 0's result - aco.hip.cpp: restore numAntsTerminated_++ in PCPU_FindOneSchedule - sched_region.hip.cpp, OptimizingScheduler.hip.cpp: guard HIP device allocation behind hipGetDeviceCount check so DEV_ACO NO runs cleanly without a GPU present - sched.ini: set DEV_ACO NO, HOST_ANTS 8; increase NO_CPU_THREADS to 8 --- example/optsched-cfg/sched.ini | 49 ++++- include/opt-sched/Scheduler/aco.h | 6 + lib/Scheduler/aco.hip.cpp | 256 +++++++++++++++--------- lib/Scheduler/bb_spill.hip.cpp | 5 +- lib/Scheduler/data_dep.hip.cpp | 2 +- lib/Scheduler/enumerator.cpp | 11 +- lib/Scheduler/random.hip.cpp | 12 +- lib/Scheduler/relaxed_sched.cpp | 5 +- lib/Scheduler/sched_region.hip.cpp | 46 +++-- lib/Wrapper/OptimizingScheduler.hip.cpp | 146 +++++++++++--- 10 files changed, 386 insertions(+), 152 deletions(-) diff --git a/example/optsched-cfg/sched.ini b/example/optsched-cfg/sched.ini index 7db698d9..b472746d 100644 --- a/example/optsched-cfg/sched.ini +++ b/example/optsched-cfg/sched.ini @@ -25,6 +25,9 @@ USE_TWO_PASS NO HEUR_ENABLED YES # ACO_ENABLED is the Ant Colony Optimization scheduler. ACO_ENABLED YES +o# MO_ENABLED toggles multiple occupancy schedule generation. +# If YES, scheduler explores multiple occupancy points; otherwise it uses a single occupancy. +MO_ENABLED NO # ENUM_ENABLED is the Branch and Bound scheduler. ENUM_ENABLED NO @@ -69,6 +72,21 @@ SECOND_PASS_LENGTH_TIMEOUT 0 # BLOCK : use the time limits in the above fields as is TIMEOUT_PER INSTR +# Whether to limit occupancy when scheduling. +# YES : apply occupancy limiting logic +# NO : do not apply occupancy limiting +SHOULD_LIMIT_OCCUPANCY NO + +# Target occupancy limit when SHOULD_LIMIT_OCCUPANCY is YES. +# 0 means unlimited. +OCCUPANCY_LIMIT 0 + +# How to determine the occupancy limit. Valid values: +# NONE : no occupancy limit source is used +# HEURISTIC : use a heuristic to decide whether to limit occupancy +# FILE : read the limit from occupancy_limits.ini +OCCUPANCY_LIMIT_SOURCE NONE + # The heuristic used for the list scheduler. Valid values are any combination of: # CP: critical path # LUC: last use count @@ -77,6 +95,8 @@ TIMEOUT_PER INSTR # NID: node ID # LLVM: LLVM’s default list scheduler order # Example: LUC_CP_NID +LIST_HEURISTIC LLVM +# Legacy key, supported for backward compatibility. HEURISTIC LLVM # The heuristic used for the enumerator. If the two pass scheduling @@ -155,6 +175,12 @@ ACO_TOURNAMENT NO #use fixd value for bias or not. If not, use ratio instaed ACO_USE_FIXED_BIAS NO +# Weighted second-pass comparison; only meaningful when USE_TWO_PASS is YES. +# VALUES: +# YES +# NO +USE_WEIGHTED_SECOND_PASS NO + #Fixed number of evaporation ACO_FIXED_BIAS 20 ACO2P_FIXED_BIAS 20 @@ -169,6 +195,24 @@ ACO2P_DECAY_FACTOR 0.2 ACO_ANT_PER_ITERATION 10 +# ACO stop iteration ranges, selected by problem size. +# These control how many stagnant iterations are allowed before ACO stops. +# If unset, defaults are: 10, 20, 50, 100 for ranges 1..4. +ACO_STOP_ITERATIONS_RANGE1 10 +ACO_STOP_ITERATIONS_RANGE2 20 +ACO_STOP_ITERATIONS_RANGE3 50 +ACO_STOP_ITERATIONS_RANGE4 100 +ACO2P_STOP_ITERATIONS_RANGE1 10 +ACO2P_STOP_ITERATIONS_RANGE2 20 +ACO2P_STOP_ITERATIONS_RANGE3 50 +ACO2P_STOP_ITERATIONS_RANGE4 100 + +# Upper bound type for the ACO ready list. +# Options: +# NO +# MIN_DEGREE +ACO_READY_LIST_UB NO + ACO_TRACE NO #If you want to use pheromone table debugging set ACO_DBG_REGIONS @@ -198,7 +242,10 @@ ACO2P_STOP_ITERATIONS 50 LLVM_MUTATIONS NO # Whether or not to run ACO on the device -DEV_ACO YES +DEV_ACO NO + +# Number of host ACO ants to use when DEV_ACO is NO +HOST_ANTS 8 # (Chris) If using the SLIL cost function, enabling this option # will force the B&B scheduler to skip DAGs with zero PERP. diff --git a/include/opt-sched/Scheduler/aco.h b/include/opt-sched/Scheduler/aco.h index b6ed43fb..cf21516a 100644 --- a/include/opt-sched/Scheduler/aco.h +++ b/include/opt-sched/Scheduler/aco.h @@ -15,6 +15,8 @@ Last Update: Jan. 2020 #include "opt-sched/Scheduler/device_vector.h" #include "llvm/ADT/ArrayRef.h" #include +#include +#include #include namespace llvm { @@ -295,6 +297,10 @@ class ACOScheduler : public ConstrainedScheduler { SchedPriorities priorities1_; SchedPriorities priorities2_; + // Pre-built per-instruction successor list (successor ptr + prdcsrNum). + // Populated before PCPU threads launch to avoid cursor race on scsrLst_. + std::vector>> instScsrs_; + }; } // namespace opt_sched diff --git a/lib/Scheduler/aco.hip.cpp b/lib/Scheduler/aco.hip.cpp index 8e9330ac..5d14113f 100644 --- a/lib/Scheduler/aco.hip.cpp +++ b/lib/Scheduler/aco.hip.cpp @@ -1,4 +1,7 @@ -#include "hip/hip_runtime.h" +#include +#include +#include + #include "opt-sched/Scheduler/aco.h" #include "opt-sched/Scheduler/config.h" #include "opt-sched/Scheduler/data_dep.h" @@ -8,13 +11,11 @@ #include "opt-sched/Scheduler/bb_spill.h" #include "opt-sched/Scheduler/dev_defines.h" // #include -#include #include "llvm/ADT/STLExtras.h" #include #include #include #include -#include #include #include @@ -55,7 +56,7 @@ double RandDouble(double min, double max) { //#endif #define RUN_PCPU 1 -#define NO_CPU_THREADS 2 +#define NO_CPU_THREADS 8 ACOScheduler::ACOScheduler(DataDepGraph *dataDepGraph, MachineModel *machineModel, InstCount upperBound, @@ -82,10 +83,10 @@ ACOScheduler::ACOScheduler(DataDepGraph *dataDepGraph, numDiffOccupancies_ = numDiffOccupancies; targetOccupancy_ = targetOccupancy; - use_dev_ACO = schedIni.GetBool("DEV_ACO"); - if(!use_dev_ACO || count_ < REGION_MIN_SIZE) - numThreads_ = schedIni.GetInt("HOST_ANTS"); - else { + use_dev_ACO = schedIni.GetBool("DEV_ACO") && dev_rgn_ != nullptr && dev_DDG_ != nullptr; + if (!use_dev_ACO || count_ < REGION_MIN_SIZE) { + numThreads_ = schedIni.GetInt("HOST_ANTS", 2); + } else { dev_rgn_->SetNumThreads(numThreads_); dev_DDG_->SetNumThreads(numThreads_); } @@ -96,7 +97,7 @@ ACOScheduler::ACOScheduler(DataDepGraph *dataDepGraph, local_decay = schedIni.GetFloat("ACO_LOCAL_DECAY"); print_aco_trace = schedIni.GetBool("ACO_TRACE"); IsTwoPassEn = schedIni.GetBool("USE_TWO_PASS"); - weightedSecondPass = schedIni.GetBool("USE_WEIGHTED_SECOND_PASS"); + weightedSecondPass = schedIni.GetBool("USE_WEIGHTED_SECOND_PASS", false); /* std::cerr << "useOldAlg===="<ScoreSum = 0; - int lastInstId = lastInst->GetNum(); + int lastInstId = lastInst ? lastInst->GetNum() : -1; // this bool is to check if stalling could be avoided bool couldAvoidStalling = false; // this bool is to check if we should currently avoid unnecessary stalls @@ -418,7 +419,7 @@ InstCount ACOScheduler::SelectInstruction(SchedInstruction *lastInst, InstCount pheromone_t MaxScore = -1; InstCount MaxScoreIndx = 0; readyLs->ScoreSum = 0; - int lastInstId = lastInst->GetNum(); + int lastInstId = lastInst ? lastInst->GetNum() : -1; // this bool is to check if stalling could be avoided bool couldAvoidStalling = false; // this bool is to check if we should currently avoid unnecessary stalls @@ -438,7 +439,7 @@ InstCount ACOScheduler::SelectInstruction(SchedInstruction *lastInst, InstCount HeurType Heur = *readyLs->getInstHeuristicAtIndex(I); pheromone_t IScore = Score(lastInstId, *readyLs->getInstIdAtIndex(I), Heur, !rgn->IsSecondPass()); if (RP0OrPositiveCount != 0 && candidateDefs > candidateLUC) - IScore = IScore * 9/10; + IScore *= 0.9; *readyLs->getInstScoreAtIndex(I) = IScore; readyLs->ScoreSum += IScore; @@ -456,14 +457,16 @@ InstCount ACOScheduler::SelectInstruction(SchedInstruction *lastInst, InstCount // add a score penalty for instructions that are not ready yet // unnecessary stalls should not be considered if current RP is low, or if we already have too many stalls if (*readyLs->getInstReadyOnAtIndex(I) > crntCycleNum_) { - if (RP0OrPositiveCount != 0) { + if (RP0OrPositiveCount != 0 || globalBestStalls_ == 0) { + // If there are RP-neutral instructions available, or we have no stall budget, + // strongly discourage waiting for this instruction. IScore = 0.0000001; } else { int cyclesNeededToWait = *readyLs->getInstReadyOnAtIndex(I) - crntCycleNum_; if (cyclesNeededToWait < globalBestStalls_) IScore = IScore * (globalBestStalls_ - cyclesNeededToWait * 2) / globalBestStalls_; - else + else IScore = IScore / globalBestStalls_; // check if any reg types used by the instructions are above the physical limit @@ -880,7 +883,11 @@ InstSchedule *ACOScheduler::FindOneSchedule(InstCount RPTarget, InstSchedule *de // initialize the aco ready list so that the start instruction is ready // The luc component is 0 since the root inst uses no instructions InstCount RootId = rootInst_->GetNum(); - HeurType RootHeuristic = kHelper1->computeKey(rootInst_, true, dataDepGraph_->RegFiles); + HeurType RootHeuristic; + if (IsSecondPass && kHelper2) + RootHeuristic = kHelper2->computeKey(rootInst_, true, 0, dataDepGraph_->RegFiles); + else + RootHeuristic = kHelper1->computeKey(rootInst_, true, dataDepGraph_->RegFiles); pheromone_t RootScore = Score(-1, RootId, RootHeuristic, !IsSecondPass); ACOReadyListEntry InitialRoot{RootId, 0, RootHeuristic, RootScore}; readyLs->addInstructionToReadyList(InitialRoot); @@ -1379,16 +1386,16 @@ FUNC_RESULT ACOScheduler::FindSchedule(InstSchedule *schedule_out, printf("target: %d\n",targetOccupancy_); if (count_ < 50) noImprovementMax = schedIni.GetInt(IsFirst ? "ACO_STOP_ITERATIONS_RANGE1" - : "ACO2P_STOP_ITERATIONS_RANGE1"); + : "ACO2P_STOP_ITERATIONS_RANGE1", 10); else if (count_ < 100) noImprovementMax = schedIni.GetInt(IsFirst ? "ACO_STOP_ITERATIONS_RANGE2" - : "ACO2P_STOP_ITERATIONS_RANGE2"); + : "ACO2P_STOP_ITERATIONS_RANGE2", 20); else if (count_ < 1000) noImprovementMax = schedIni.GetInt(IsFirst ? "ACO_STOP_ITERATIONS_RANGE3" - : "ACO2P_STOP_ITERATIONS_RANGE3"); + : "ACO2P_STOP_ITERATIONS_RANGE3", 50); else noImprovementMax = schedIni.GetInt(IsFirst ? "ACO_STOP_ITERATIONS_RANGE4" - : "ACO2P_STOP_ITERATIONS_RANGE4"); + : "ACO2P_STOP_ITERATIONS_RANGE4", 100); if (dev_AcoSchdulr) dev_AcoSchdulr->noImprovementMax = noImprovementMax; @@ -1427,7 +1434,6 @@ FUNC_RESULT ACOScheduler::FindSchedule(InstSchedule *schedule_out, for (int i = 0; i < pheromone_size; i++) pheromone_[i] = initialValue_; #endif - std::cerr << "initialValue_" << initialValue_ << std::endl; InstSchedule *bestSchedule = InitialSchedule; // check if heuristic schedule is better than the initial @@ -1975,6 +1981,13 @@ void ACOScheduler::UpdatePheromone(InstSchedule *schedule, bool isIterationBest, PrintPheromone(blockOccupancyNum); #else // host version of function + // Defensive checks: ensure schedule is valid and ScRelMax non-zero + if (!schedule) + return; + + if (ScRelMax == 0) + ScRelMax = 1; // avoid divide-by-zero + // I wish InstSchedule allowed you to just iterate over it, but it's got this // cycle and slot thing which needs to be accounted for InstCount instNum, cycleNum, slotNum; @@ -1985,8 +1998,10 @@ void ACOScheduler::UpdatePheromone(InstSchedule *schedule, bool isIterationBest, pheromone_t deposition = fmax((1 - portion) * MAX_DEPOSITION_MINUS_MIN, 0) + MIN_DEPOSITION; pheromone_t *pheromone; - while (instNum != INVALID_VALUE) { + while (instNum != INVALID_VALUE) { SchedInstruction *inst = dataDepGraph_->GetInstByIndx(instNum); + if (!inst) + break; pheromone = &Pheromone(lastInst, inst); #if USE_ACS @@ -2004,11 +2019,13 @@ void ACOScheduler::UpdatePheromone(InstSchedule *schedule, bool isIterationBest, schedule->ResetInstIter(); #if !USE_ACS - // decay pheromone - for (int i = 0; i < count_; i++) { - for (int j = 0; j < count_; j++) { - pheromone = &Pheromone(i, j); - *pheromone *= (1 - decay_factor); + // decay pheromone (only if count_ is sane) + if (count_ > 0) { + for (int i = 0; i < count_; i++) { + for (int j = 0; j < count_; j++) { + pheromone = &Pheromone(i, j); + *pheromone *= (1 - decay_factor); + } } } #endif @@ -2187,14 +2204,19 @@ inline void ACOScheduler::UpdateACOReadyList(SchedInstruction *inst, bool IsSeco if (wasLastPrdcsr) { // If all other predecessors of this successor have been scheduled then // we now know in which cycle this successor will become ready. - HeurType HeurWOLuc = kHelper1->computeKey(crntScsr, false, dataDepGraph_->RegFiles); + HeurType HeurWOLuc; + if (IsSecondPass && kHelper2) + HeurWOLuc = kHelper2->computeKey(crntScsr, false, heurChoice, dataDepGraph_->RegFiles); + else + HeurWOLuc = kHelper1->computeKey(crntScsr, false, dataDepGraph_->RegFiles); readyLs->addInstructionToReadyList(ACOReadyListEntry{crntScsr->GetNum(), scsrRdyCycle, HeurWOLuc, 0}); } } // Make sure the scores are valid. The scheduling of an instruction may // have increased another instruction's LUC Score - PriorityEntry LUCEntry = kHelper1->getPriorityEntry(LSH_LUC); + PriorityEntry LUCEntry = (IsSecondPass && kHelper2) ? kHelper2->getPriorityEntry(LSH_LUC, 1) + : kHelper1->getPriorityEntry(LSH_LUC); RP0OrPositiveCount = 0; for (InstCount I = 0; I < readyLs->getReadyListSize(); ++I) { //we first get the heuristic without the LUC component, add the LUC @@ -2207,17 +2229,14 @@ inline void ACOScheduler::UpdateACOReadyList(SchedInstruction *inst, bool IsSeco LUCVal <<= LUCEntry.Offset; Heur &= LUCVal; } - if (RP0OrPositiveCount) { - if (*dev_readyLs->getInstReadyOnAtIndex(I) > crntCycleNum_) - continue; + if (*readyLs->getInstReadyOnAtIndex(I) > crntCycleNum_) + continue; - SchedInstruction *candidateInst = dataDepGraph_->GetInstByIndx(CandidateId); - HeurType candidateLUC = candidateInst->GetLastUseCnt(); - int16_t candidateDefs = candidateInst->GetDefCnt(); - if (candidateDefs <= candidateLUC) { - RP0OrPositiveCount = RP0OrPositiveCount + 1; - } - } + SchedInstruction *candidateInst = dataDepGraph_->GetInstByIndx(CandidateId); + HeurType candidateLUC = candidateInst->GetLastUseCnt(); + int16_t candidateDefs = candidateInst->GetDefCnt(); + if (candidateDefs <= candidateLUC) + RP0OrPositiveCount = RP0OrPositiveCount + 1; } #endif } @@ -2432,6 +2451,19 @@ InstSchedule *ACOScheduler::FindManyCPUSchedule(InstCount RPTarget) { //allocate new scheduler variables for parallel PCPUACOSchedVars *pcpu_sched_vars = AllocPCPUACOSchedVars(NO_CPU_THREADS); + // Pre-build per-instruction successor lists so parallel threads can iterate + // successors without racing on scsrLst_->rtrvEntry_ (internal cursor state). + instScsrs_.assign(count_, {}); + for (int i = 0; i < count_; i++) { + SchedInstruction *si = dataDepGraph_->GetInstByIndx(i); + if (!si) + continue; + InstCount prdcsrNum; + for (SchedInstruction *scsr = si->GetFrstScsr(&prdcsrNum); + scsr != NULL; scsr = si->GetNxtScsr(&prdcsrNum)) + instScsrs_[i].emplace_back(scsr, prdcsrNum); + } + std::vector PCPUThreads; InstSchedule **cpuScheds = new InstSchedule*[NO_CPU_THREADS](); ((BBWithSpill*)rgn_)->AllocParallelCPUVars(NO_CPU_THREADS); @@ -2449,22 +2481,31 @@ InstSchedule *ACOScheduler::FindManyCPUSchedule(InstCount RPTarget) { } ((BBWithSpill*)rgn_)->FreeParallelCPUVars(NO_CPU_THREADS); - //logic to find best schedule should be added here - InstSchedule *result = cpuScheds[0]; - MaxPriorityInv = pcpu_sched_vars[0].MaxPriorityInv; - //the index with the best schedule should also grab the pcpu_sched_vars[i].MaxPriorityInv - //and update the global MaxPriorityInv + // Choose the best schedule from the CPU threads. + InstSchedule *result = nullptr; + MaxPriorityInv = 0; - for (int i = 1; i < NO_CPU_THREADS; i++) { - delete cpuScheds[i]; + for (int i = 0; i < NO_CPU_THREADS; i++) { + if (!cpuScheds[i]) + continue; + + if (!result || shouldReplaceSchedule(result, cpuScheds[i], false, RPTarget, targetOccupancy_)) { + if (result && result != cpuScheds[i]) + delete result; + result = cpuScheds[i]; + MaxPriorityInv = pcpu_sched_vars[i].MaxPriorityInv; + } else { + delete cpuScheds[i]; + } } + //free new scheduler variables for parallel FreePCPUACOSchedVars(pcpu_sched_vars, NO_CPU_THREADS); pcpu_sched_vars = nullptr; delete[] cpuScheds; PCPU_FreeSchedInsts(NO_CPU_THREADS); - printf("Ending FindManyCPUSchedule"); + Logger::Info("FindManyCPUSchedule done"); return result; } @@ -2594,8 +2635,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, // schedule construction if (((BBWithSpill*)rgn_)->PCPU_GetCrntSpillCost(pcpu) > RPTarget) { // end schedule construction - // keep track of ants terminated - //numAntsTerminated_++; + numAntsTerminated_++; pcpu_sched_vars.readyLs->clearReadyList(); delete schedule; Logger::Info("Thread %d, returned early", thread); @@ -2637,35 +2677,54 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC // calculate MaxScoringInst, and ScoreSum pheromone_t MaxScore = -1; InstCount MaxScoreIndx = 0; - pcpu_sched_vars.readyLs->ScoreSum = 0; - int lastInstId = lastInst->GetNum(); + ACOReadyList *readyLs = pcpu_sched_vars.readyLs; + if (!readyLs) + return -1; + + size_t readyListSize = readyLs->getReadyListSize(); + if (readyListSize == 0) + return -1; + + readyLs->ScoreSum = 0; + int lastInstId = lastInst ? lastInst->GetNum() : -1; // this bool is to check if stalling could be avoided bool couldAvoidStalling = false; // this bool is to check if we should currently avoid unnecessary stalls // because RP is low or we have too many stalls in the schedule bool RPIsHigh = false; bool tooManyStalls = totalStalls >= globalBestStalls_ * 5 / 10; - pcpu_sched_vars.readyLs->ScoreSum = 0; - for (InstCount I = 0; I < pcpu_sched_vars.readyLs->getReadyListSize(); ++I) { + for (InstCount I = 0; I < readyListSize; ++I) { RPIsHigh = false; - InstCount CandidateId = *pcpu_sched_vars.readyLs->getInstIdAtIndex(I); + // Fetch pointers once and validate them to avoid crashes from null returns + InstCount *candidateIdPtr = readyLs->getInstIdAtIndex(I); + HeurType *heurPtr = pcpu_sched_vars.readyLs->getInstHeuristicAtIndex(I); + pheromone_t *scorePtr = pcpu_sched_vars.readyLs->getInstScoreAtIndex(I); + InstCount *readyOnPtr = pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I); + + if (!candidateIdPtr || !heurPtr || !scorePtr || !readyOnPtr) + continue; + + InstCount CandidateId = *candidateIdPtr; + if (CandidateId >= dataDepGraph_->GetInstCnt()) + continue; SchedInstruction *candidateInst = dataDepGraph_->GetInstByIndx(CandidateId); + if (!candidateInst) + continue; HeurType candidateLUC = candidateInst->PCPU_GetLastUseCnt(thread); int16_t candidateDefs = candidateInst->GetDefCnt(); // compute the score - HeurType Heur = *pcpu_sched_vars.readyLs->getInstHeuristicAtIndex(I); - pheromone_t IScore = PCPU_Score(lastInstId, *pcpu_sched_vars.readyLs->getInstIdAtIndex(I), Heur, !rgn->IsSecondPass(), pcpu_sched_vars); + HeurType Heur = *heurPtr; + pheromone_t IScore = PCPU_Score(lastInstId, CandidateId, Heur, !rgn->IsSecondPass(), pcpu_sched_vars); if (pcpu_sched_vars.RP0OrPositiveCount != 0 && candidateDefs > candidateLUC) - IScore = IScore * 9/10; + IScore *= 0.9; - *pcpu_sched_vars.readyLs->getInstScoreAtIndex(I) = IScore; - pcpu_sched_vars.readyLs->ScoreSum += IScore; + *scorePtr = IScore; if (currentlyWaiting) { - // if currently waiting on an instruction, do not consider semi-ready instructions - if (*pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) > pcpu_sched_vars.crntCycleNum) + // if currently waiting on an instruction, do not consider semi-ready instructions + if (*readyOnPtr > pcpu_sched_vars.crntCycleNum) continue; // as well as instructions with a net negative impact on RP @@ -2675,28 +2734,38 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC // add a score penalty for instructions that are not ready yet // unnecessary stalls should not be considered if current RP is low, or if we already have too many stalls - if (*pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) > pcpu_sched_vars.crntCycleNum) { - if (pcpu_sched_vars.RP0OrPositiveCount != 0) { + if (*readyOnPtr > pcpu_sched_vars.crntCycleNum) { + if (pcpu_sched_vars.RP0OrPositiveCount != 0 || globalBestStalls_ == 0) { + // If there are RP-neutral instructions available, or we have no stall budget, + // strongly discourage waiting for this instruction. IScore = 0.0000001; } else { - int cyclesNeededToWait = *pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) - pcpu_sched_vars.crntCycleNum; + int cyclesNeededToWait = *readyOnPtr - pcpu_sched_vars.crntCycleNum; if (cyclesNeededToWait < globalBestStalls_) IScore = IScore * (globalBestStalls_ - cyclesNeededToWait * 2) / globalBestStalls_; - else + else IScore = IScore / globalBestStalls_; // check if any reg types used by the instructions are above the physical limit SchedInstruction *tempInst = dataDepGraph_->GetInstByIndx(*pcpu_sched_vars.readyLs->getInstIdAtIndex(I)); + if (!tempInst) + continue; RegIndxTuple *uses; Register *use; uint16_t usesCount = tempInst->GetUses(uses); - for (uint16_t i = 0; i < usesCount; i++) { - use = dataDepGraph_->getRegByTuple(&uses[i]); - int16_t regType = use->GetType(); - if ( ((BBWithSpill *)rgn)->PCPU_IsRPHigh(regType, pcpu) ) { - RPIsHigh = true; - break; + if (usesCount == 0 || !uses) + ; // nothing to check + else { + for (uint16_t i = 0; i < usesCount; i++) { + use = dataDepGraph_->getRegByTuple(&uses[i]); + if (!use) + continue; + int16_t regType = use->GetType(); + if ( ((BBWithSpill *)rgn)->PCPU_IsRPHigh(regType, pcpu) ) { + RPIsHigh = true; + break; + } } } @@ -2718,13 +2787,16 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC IScore = 0.0000001; *pcpu_sched_vars.readyLs->getInstScoreAtIndex(I) = IScore; pcpu_sched_vars.readyLs->ScoreSum += IScore; - - if(IScore > MaxScore) { + + if (IScore > MaxScore) { MaxScoreIndx = I; MaxScore = IScore; } } + if (MaxScore < 0) + return -1; + //generate the random numbers that we will need for deciding if //we are going to use the fixed bias or if we are going to use //fitness proportional selection. Generate the number used for @@ -2746,20 +2818,24 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC //indices of the max and fp choice instructions //The only branch in this code is the branch for deciding to stay in the loop vs exit the loop //this will diverge if two ants ready lists are of different sizes - // select the instruction index for fp choice - size_t fpIndx = 0; - for (size_t i = 0; i < pcpu_sched_vars.readyLs->getReadyListSize(); ++i) { + // select the instruction index for fp choice + size_t fpIndx = 0; + for (size_t i = 0; i < readyListSize; ++i) { point -= *pcpu_sched_vars.readyLs->getInstScoreAtIndex(i); if (point <= 0) { fpIndx = i; break; } } + if (point > 0 && readyListSize > 0) + fpIndx = readyListSize - 1; - //finally we pick whether we will return the fp choice or max score inst w/o using a branch + // finally we pick whether we will return the fp choice or max score inst w/o using a branch size_t indx; bool UseMax = (rand < choose_best_chance) || currentlyWaiting; indx = UseMax ? MaxScoreIndx : fpIndx; + if (indx >= readyListSize) + indx = readyListSize - 1; if (couldAvoidStalling && *pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(indx) > pcpu_sched_vars.crntCycleNum) unnecessarilyStalling = true; @@ -2775,11 +2851,14 @@ inline void ACOScheduler::PCPU_UpdateACOReadyList(SchedInstruction *inst, bool I int thread, PCPUACOSchedVars &pcpu_sched_vars, int heurChoice){ - InstCount prdcsrNum, scsrRdyCycle; - //Logger::Info("Start UpdateACOReadyList, %d", thread); + if (!pcpu_sched_vars.readyLs) + return; + InstCount scsrRdyCycle; + InstCount instIdx = inst->GetNum(); + //Logger::Info("Start UpdateACOReadyList, %d", thread); // Notify each successor of this instruction that it has been scheduled. - for (SchedInstruction *crntScsr = inst->GetFrstScsr(&prdcsrNum); - crntScsr != NULL; crntScsr = inst->GetNxtScsr(&prdcsrNum)) { + // Use pre-built instScsrs_ to avoid cursor race on scsrLst_->rtrvEntry_. + for (auto &[crntScsr, prdcsrNum] : instScsrs_[instIdx]) { bool wasLastPrdcsr = crntScsr->PCPU_PrdcsrSchduld(prdcsrNum, pcpu_sched_vars.crntCycleNum, scsrRdyCycle, thread); @@ -2806,17 +2885,14 @@ inline void ACOScheduler::PCPU_UpdateACOReadyList(SchedInstruction *inst, bool I LUCVal <<= LUCEntry.Offset; Heur &= LUCVal; } - if (pcpu_sched_vars.RP0OrPositiveCount) { - if (*pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) > pcpu_sched_vars.crntCycleNum) - continue; + if (*pcpu_sched_vars.readyLs->getInstReadyOnAtIndex(I) > pcpu_sched_vars.crntCycleNum) + continue; - SchedInstruction *candidateInst = dataDepGraph_->GetInstByIndx(CandidateId); - HeurType candidateLUC = candidateInst->PCPU_GetLastUseCnt(thread); - int16_t candidateDefs = candidateInst->GetDefCnt(); - if (candidateDefs <= candidateLUC) { - pcpu_sched_vars.RP0OrPositiveCount = pcpu_sched_vars.RP0OrPositiveCount + 1; - } - } + SchedInstruction *candidateInst = dataDepGraph_->GetInstByIndx(CandidateId); + HeurType candidateLUC = candidateInst->PCPU_GetLastUseCnt(thread); + int16_t candidateDefs = candidateInst->GetDefCnt(); + if (candidateDefs <= candidateLUC) + pcpu_sched_vars.RP0OrPositiveCount = pcpu_sched_vars.RP0OrPositiveCount + 1; } //Logger::Info("Finish UpdateACOReadyList, %d", thread); } diff --git a/lib/Scheduler/bb_spill.hip.cpp b/lib/Scheduler/bb_spill.hip.cpp index 05d5c559..2c51b9c6 100644 --- a/lib/Scheduler/bb_spill.hip.cpp +++ b/lib/Scheduler/bb_spill.hip.cpp @@ -1,3 +1,5 @@ +#include + #include "opt-sched/Scheduler/bb_spill.h" #include "opt-sched/Scheduler/config.h" #include "opt-sched/Scheduler/data_dep.h" @@ -20,7 +22,6 @@ #include #include #include -#include // #define IS_DEBUG_REG_PRESSURE 1 extern bool OPTSCHED_gPrintSpills; @@ -2117,4 +2118,4 @@ InstCount BBWithSpill::PCPU_CmputCost_(InstSchedule *sched, COST_COMP_MODE compM ParallelCPUVars &BBWithSpill::GetPCPUVars(int thread){ return pcpu_vars_[thread]; -} \ No newline at end of file +} diff --git a/lib/Scheduler/data_dep.hip.cpp b/lib/Scheduler/data_dep.hip.cpp index b63ba2a6..929c4438 100644 --- a/lib/Scheduler/data_dep.hip.cpp +++ b/lib/Scheduler/data_dep.hip.cpp @@ -1,3 +1,4 @@ +#include #include #include #include @@ -16,7 +17,6 @@ #include "opt-sched/Scheduler/config.h" #include "llvm/ADT/STLExtras.h" #include "llvm/Support/Debug.h" -#include // only print pressure if enabled by sched.ini extern bool OPTSCHED_gPrintSpills; diff --git a/lib/Scheduler/enumerator.cpp b/lib/Scheduler/enumerator.cpp index 313411d8..cbc7f7b6 100644 --- a/lib/Scheduler/enumerator.cpp +++ b/lib/Scheduler/enumerator.cpp @@ -1,3 +1,9 @@ +#include +#include +#include +#include +#include + #include "opt-sched/Scheduler/enumerator.h" #include "opt-sched/Scheduler/bb_spill.h" #include "opt-sched/Scheduler/hist_table.h" @@ -5,11 +11,6 @@ #include "opt-sched/Scheduler/random.h" #include "opt-sched/Scheduler/stats.h" #include "opt-sched/Scheduler/utilities.h" -#include -#include -#include -#include -#include using namespace llvm::opt_sched; diff --git a/lib/Scheduler/random.hip.cpp b/lib/Scheduler/random.hip.cpp index 4c6d7228..2d6348b4 100644 --- a/lib/Scheduler/random.hip.cpp +++ b/lib/Scheduler/random.hip.cpp @@ -47,7 +47,7 @@ static uint32_t y[] = { static uint32_t randNum; void GenerateNextNumber() { - std::scoped_lock{m}; + std::scoped_lock lock(m); randNum = y[j] + y[k]; y[k] = randNum; if (--j < 0) @@ -58,7 +58,7 @@ void GenerateNextNumber() { } void RandomGen::SetSeed(int32_t iseed) { - std::scoped_lock{m}; + std::scoped_lock lock(m); j = 23; k = 54; @@ -75,19 +75,19 @@ void RandomGen::SetSeed(int32_t iseed) { } uint32_t RandomGen::GetRand32WithinRange(uint32_t min, uint32_t max) { - std::scoped_lock{m}; + std::scoped_lock lock(m); GenerateNextNumber(); return randNum % (max - min + 1) + min; } uint32_t RandomGen::GetRand32() { - std::scoped_lock{m}; + std::scoped_lock lock(m); GenerateNextNumber(); return randNum; } uint64_t RandomGen::GetRand64() { - std::scoped_lock{m}; + std::scoped_lock lock(m); uint64_t rand64; GenerateNextNumber(); @@ -101,7 +101,7 @@ uint64_t RandomGen::GetRand64() { } void RandomGen::GetRandBits(uint16_t bitCnt, unsigned char *dest) { - std::scoped_lock{m}; + std::scoped_lock lock(m); uint16_t bytesNeeded = (bitCnt + 7) / 8; uint16_t index = 0; diff --git a/lib/Scheduler/relaxed_sched.cpp b/lib/Scheduler/relaxed_sched.cpp index 9df3cd92..0611ce76 100644 --- a/lib/Scheduler/relaxed_sched.cpp +++ b/lib/Scheduler/relaxed_sched.cpp @@ -1,8 +1,9 @@ +#include +#include + #include "opt-sched/Scheduler/relaxed_sched.h" #include "opt-sched/Scheduler/logger.h" #include "opt-sched/Scheduler/utilities.h" -#include -#include using namespace llvm::opt_sched; diff --git a/lib/Scheduler/sched_region.hip.cpp b/lib/Scheduler/sched_region.hip.cpp index 3a44f5d9..65f00050 100644 --- a/lib/Scheduler/sched_region.hip.cpp +++ b/lib/Scheduler/sched_region.hip.cpp @@ -1,4 +1,6 @@ -#include "hip/hip_runtime.h" +#include +#include +#include #include #include #include @@ -18,11 +20,9 @@ #include "opt-sched/Scheduler/stats.h" #include "opt-sched/Scheduler/utilities.h" #include "opt-sched/Scheduler/dev_defines.h" -#include #include "llvm/ADT/SmallString.h" #include "llvm/Support/ErrorHandling.h" #include "llvm/Support/FileSystem.h" -#include extern bool OPTSCHED_gPrintSpills; @@ -254,21 +254,25 @@ FUNC_RESULT SchedRegion::FindOptimalSchedule( Config &schedIni = SchedulerOptions::getInstance(); bool HeuristicSchedulerEnabled = schedIni.GetBool("HEUR_ENABLED"); bool AcoSchedulerEnabled = schedIni.GetBool("ACO_ENABLED"); - multipleOccupancies = schedIni.GetBool("MO_ENABLED"); - printf("multipleOccupancies is %d",multipleOccupancies); + multipleOccupancies = schedIni.GetBool("MO_ENABLED", false); + printf("multipleOccupancies is %d", multipleOccupancies); bool BbSchedulerEnabled = isBbEnabled(schedIni, rgnTimeout); - unsigned long randSeed = (unsigned long) schedIni.GetInt("RANDOM_SEED"); - bool devACOEnabled = schedIni.GetBool("DEV_ACO"); + unsigned long randSeed = (unsigned long) schedIni.GetInt("RANDOM_SEED", 0); + bool devACOEnabled = schedIni.GetBool("DEV_ACO", false); int numBlocks; - if (devACOEnabled && dataDepGraph_->GetInstCnt() >= REGION_MIN_SIZE) - numBlocks = schedIni.GetBool("ACO_MANY_ANTS_ENABLED") && dataDepGraph_->GetInstCnt() > MANY_ANT_MIN_SIZE ? - schedIni.GetInt("ACO_MANY_ANTS_PER_ITERATION_BLOCKS") : schedIni.GetInt("ACO_DEVICE_ANT_PER_ITERATION_BLOCKS"); - else - numBlocks = schedIni.GetInt("HOST_ANTS"); + if (devACOEnabled && dataDepGraph_->GetInstCnt() >= REGION_MIN_SIZE) { + bool manyAntsEnabled = schedIni.GetBool("ACO_MANY_ANTS_ENABLED", false); + int manyAntBlocks = static_cast(schedIni.GetInt("ACO_MANY_ANTS_PER_ITERATION_BLOCKS", 1)); + int deviceAntBlocks = static_cast(schedIni.GetInt("ACO_DEVICE_ANT_PER_ITERATION_BLOCKS", 1)); + numBlocks = manyAntsEnabled && dataDepGraph_->GetInstCnt() > MANY_ANT_MIN_SIZE ? + manyAntBlocks : deviceAntBlocks; + } else { + numBlocks = static_cast(schedIni.GetInt("HOST_ANTS", 2)); + } if (AcoSchedulerEnabled) { - AcoBeforeEnum = schedIni.GetBool("ACO_BEFORE_ENUM"); - AcoAfterEnum = schedIni.GetBool("ACO_AFTER_ENUM"); + AcoBeforeEnum = schedIni.GetBool("ACO_BEFORE_ENUM", false); + AcoAfterEnum = schedIni.GetBool("ACO_AFTER_ENUM", false); } if (!HeuristicSchedulerEnabled && !AcoBeforeEnum) { @@ -388,7 +392,7 @@ FUNC_RESULT SchedRegion::FindOptimalSchedule( CmputNormCost_(lstSched, CCM_DYNMC, hurstcExecCost, true); hurstcCost_ = lstSched->GetCost(); InstCount maxIndependentInstructions = 0; - std::string readyListUB = schedIni.GetString("ACO_READY_LIST_UB"); + std::string readyListUB = schedIni.GetString("ACO_READY_LIST_UB", "NO"); if (readyListUB == "NO" || readyListUB == "MIN_DEGREE") { for (int i = 0; i < dataDepGraph_->GetInstCnt(); i++) { int independentInstructions = dataDepGraph_->GetInstCnt() - dataDepGraph_->GetInstByIndx(i)->GetRcrsvPrdcsrCnt() - dataDepGraph_->GetInstByIndx(i)->GetRcrsvScsrCnt(); @@ -1143,6 +1147,18 @@ FUNC_RESULT SchedRegion::runACO(InstSchedule *ReturnSched, // to fit in device memory Logger::Info("This DDG has %d edges", dataDepGraph_->GetEdgeCnt()); Logger::Info("CP Distance: %d", dataDepGraph_->GetRootInst()->GetCrntLwrBound(DIR_BKWRD) + 1); + if (devACOEnabled) { + int hipDeviceCount = 0; + hipError_t hipErr = hipGetDeviceCount(&hipDeviceCount); + if (hipErr != hipSuccess || hipDeviceCount == 0) { + Logger::Info("DEV_ACO disabled in runACO: no HIP devices found (%s)", + hipGetErrorString(hipErr)); + devACOEnabled = false; + } else { + gpuErrchk(hipSetDevice(0)); + } + } + if (devACOEnabled && dataDepGraph_->GetInstCnt() >= REGION_MIN_SIZE) { // Allocate and Copy data to device for parallel ACO size_t memSize; diff --git a/lib/Wrapper/OptimizingScheduler.hip.cpp b/lib/Wrapper/OptimizingScheduler.hip.cpp index 8dd2c4b7..acd159b9 100644 --- a/lib/Wrapper/OptimizingScheduler.hip.cpp +++ b/lib/Wrapper/OptimizingScheduler.hip.cpp @@ -43,6 +43,7 @@ #include #include #include +#include #define DEBUG_TYPE "optsched" @@ -283,27 +284,77 @@ void ScheduleDAGOptSched::initSchedulers() { } void ScheduleEvaluator::recordSchedule(int schedIndex) { - schedCount++; - if (schedIndex >= storedSchedules.size()) - storedSchedules.resize(schedIndex + 1); + if (schedIndex < 0) + return; + + if (!DAG.BB || DAG.RegionBegin == DAG.RegionEnd) + return; + + dbgs() << "[OPTSCHED_DEBUG] recordSchedule idx=" << schedIndex + << " this=" << (const void *)this + << " DAG=" << (const void *)&DAG + << " DAG.BB=" << (const void *)DAG.BB + << " RegionBegin==RegionEnd=" << (DAG.RegionBegin == DAG.RegionEnd) + << " storedSchedules.size()=" << storedSchedules.size() << "\n"; + + if ((size_t)schedIndex >= storedSchedules.size()) + storedSchedules.resize((size_t)schedIndex + 1); + + dbgs() << "[OPTSCHED_DEBUG] recordSchedule idx=" << schedIndex + << " DAG.BB=" << (void *)DAG.BB + << " RegionBegin==RegionEnd=" << (DAG.RegionBegin == DAG.RegionEnd) + << " storedSchedules.size()=" << storedSchedules.size() << "\n"; storedSchedules[schedIndex].clear(); - storedSchedules[schedIndex].reserve(DAG.NumRegionInstrs); + schedCount++; - for (auto &MI : DAG) - storedSchedules[schedIndex].push_back(&MI); + const char *dbgEnv = std::getenv("OPTSCHED_DEBUG"); + const bool doDebug = dbgEnv && std::atoi(dbgEnv) != 0; + if (doDebug) { + dbgs() << "[OPTSCHED_DEBUG] recordSchedule idx=" << schedIndex + << " DAG.BB=" << (void *)DAG.BB + << " RegionBegin==RegionEnd=" << (DAG.RegionBegin == DAG.RegionEnd) + << " storedSchedules.size()=" << storedSchedules.size() << "\n"; + } + + // Safely iterate over scheduled instructions in the region + if (!DAG.BB || DAG.RegionBegin == DAG.RegionEnd) + return; + + for (auto I = DAG.RegionBegin; I != DAG.RegionEnd; ++I) { + if (!I->isDebugInstr()) { + storedSchedules[schedIndex].push_back(&*I); + } + } + + if (doDebug) { + dbgs() << "[OPTSCHED_DEBUG] storedSchedules[" << schedIndex << "] size=" + << storedSchedules[schedIndex].size() << "\n"; + int __mi_i = 0; + for (MachineInstr *MI : storedSchedules[schedIndex]) { + dbgs() << " [" << __mi_i++ << "] "; + MI->print(dbgs()); + dbgs() << '\n'; + } + } } // revert to a previous schedule // schedIndex 0 is the previous schedule, 1 is the schedule at highest occupancy void ScheduleEvaluator::revertScheduling(int schedIndex) { + if (schedIndex < 0 || schedCount == 0 || storedSchedules.empty()) + return; + + if (schedIndex >= schedCount) + schedIndex = schedCount - 1; + + if ((size_t)schedIndex >= storedSchedules.size() || storedSchedules[schedIndex].empty()) + return; + DAG.RegionEnd = DAG.RegionBegin; int SkippedDebugInstr = 0; // Logger::Info("Reverting Scheduling Number of Scheds: %d, index: %d, size is %d", schedCount, schedIndex, storedSchedules.size()); - if (schedIndex >= schedCount) { - schedIndex = schedCount - 1; - } for (MachineInstr *MI : storedSchedules[schedIndex]) { if (MI->isDebugInstr()) { @@ -382,28 +433,33 @@ void ScheduleEvaluator::calculateRPAfter(int schedIndex) { unsigned ScheduleEvaluator::getOccAtIndex(int schedIndex) const { const GCNSubtarget &ST = DAG.MF.getSubtarget(); - // printf("storedSchedules.size() = %d | ",storedSchedules.size()); - if (schedIndex >= storedSchedules.size()) { - // printf("schedIndex: %d schedIndex too high\n", schedIndex); + if (schedIndex < 0 || (size_t)schedIndex >= storedSchedules.size()) { + if (SchedRP.empty()) + return 0; return SchedRP.back().getOccupancy(ST); } - // printf("schedIndex: %d\n", schedIndex); return SchedRP[schedIndex].getOccupancy(ST); } int ScheduleEvaluator::getSchedLength(int schedIndex) { + if (schedIndex < 0 || (size_t)schedIndex >= storedSchedules.size()) + return 0; return storedSchedules[schedIndex].size(); } int64_t ScheduleEvaluator::getILPAtIndex(int schedIndex) const { - if (schedIndex >= storedSchedules.size()) { + if (schedIndex < 0 || (size_t)schedIndex >= storedSchedules.size()) { + if (SchedILP.empty()) + return 0; return SchedILP.back(); } return SchedILP[schedIndex]; } int64_t ScheduleEvaluator::getWeightedILPAtIndex(int schedIndex) const { - if (schedIndex >= storedSchedules.size()) { + if (schedIndex < 0 || (size_t)schedIndex >= storedSchedules.size()) { + if (SchedILP.empty()) + return 0; return SchedILP.back() * Frequency; } return SchedILP[schedIndex] * Frequency; @@ -556,11 +612,19 @@ void ScheduleDAGOptSched::schedule() { std::string(":") + std::to_string(RegionNumber); + // Ensure we have a ScheduleEvaluator for this region. + LLVM_DEBUG(dbgs() << "[OPTSCHED_DEBUG] schedule entry this=" << (const void *)this + << " RegionNumber=" << RegionNumber + << " BB=" << (const void *)BB + << " RegionBegin==RegionEnd=" << (RegionBegin == RegionEnd) + << " SchedEvals.size=" << SchedEvals.size() << "\n"); + while ((size_t)RegionNumber >= SchedEvals.size()) + SchedEvals.emplace_back(*this); + // If two pass scheduling is enabled then // first just record the scheduling region. if (OptSchedEnabled && TwoPassEnabled && !TwoPassSchedulingStarted) { Regions.push_back(std::make_pair(RegionBegin, RegionEnd)); - SchedEvals.emplace_back(*this); LLVM_DEBUG( dbgs() << "Recording scheduling region before scheduling with two pass " "scheduler...\n"); @@ -680,6 +744,12 @@ void ScheduleDAGOptSched::schedule() { // Revord MachineInstr order in the first pass for a possible revert if // scheduling makes things worse. if (!SecondPass) { + LLVM_DEBUG(dbgs() << "[OPTSCHED_DEBUG] before first recordSchedule this=" + << (const void *)this + << " RegionNumber=" << RegionNumber + << " BB=" << (const void *)BB + << " RegionBegin==RegionEnd=" << (RegionBegin == RegionEnd) + << "\n"); SchedEval.recordSchedule(0); SchedEval.calculateRPBefore(); SchedEval.calcualteILPBefore(); @@ -713,16 +783,25 @@ void ScheduleDAGOptSched::schedule() { // Prepare for device scheduling by increasing heap size and copying machMdl bool dev_ACOEnabled = schedIni.GetBool("DEV_ACO"); if (dev_ACOEnabled && dev_MM == NULL && NumRegionInstrs + 2 >= REGION_MIN_SIZE) { - // Copy MachineModel to device for use during DevListSched. - // Allocate device memory - gpuErrchk(hipMallocManaged((void**)&dev_MM, sizeof(MachineModel))); - // Copy machMdl_ to device - gpuErrchk(hipMemcpy(dev_MM, MM.get(), sizeof(MachineModel), - hipMemcpyHostToDevice)); - // Copy over all pointers to device - MM.get()->CopyPointersToDevice(dev_MM); - // make sure mallocmanaged mem is copied to device before kernel start - gpuErrchk(hipMemPrefetchAsync(dev_MM, sizeof(MachineModel), 0)); + int hipDeviceCount = 0; + hipError_t hipErr = hipGetDeviceCount(&hipDeviceCount); + if (hipErr != hipSuccess || hipDeviceCount == 0) { + LLVM_DEBUG(dbgs() << "[OPTSCHED_DEBUG] DEV_ACO disabled: no HIP devices found (" + << hipGetErrorString(hipErr) << ")\n"); + dev_ACOEnabled = false; + } else { + gpuErrchk(hipSetDevice(0)); + // Copy MachineModel to device for use during DevListSched. + // Allocate device memory + gpuErrchk(hipMallocManaged((void**)&dev_MM, sizeof(MachineModel))); + // Copy machMdl_ to device + gpuErrchk(hipMemcpy(dev_MM, MM.get(), sizeof(MachineModel), + hipMemcpyHostToDevice)); + // Copy over all pointers to device + MM.get()->CopyPointersToDevice(dev_MM); + // make sure mallocmanaged mem is copied to device before kernel start + gpuErrchk(hipMemPrefetchAsync(dev_MM, sizeof(MachineModel), 0)); + } } // create region @@ -843,6 +922,11 @@ void ScheduleDAGOptSched::schedule() { } } placeDebugValues(); + LLVM_DEBUG(dbgs() << "[OPTSCHED_DEBUG] before recordSchedule idx=" << schedIndex + << " this=" << (const void *)this + << " BB=" << (const void *)BB + << " RegionBegin==RegionEnd=" << (RegionBegin == RegionEnd) + << "\n"); SchedEval.recordSchedule(schedIndex); SchedEval.calculateRPAfter(schedIndex); SchedEval.calculateILPAfter(schedIndex); @@ -988,7 +1072,8 @@ void ScheduleDAGOptSched::loadOptSchedConfig() { EnumStalls = schedIni.GetBool("ENUMERATE_STALLS"); SCW = schedIni.GetInt("SPILL_COST_WEIGHT"); LowerBoundAlgorithm = parseLowerBoundAlgorithm(); - HeuristicPriorities = parseHeuristic(schedIni.GetString("LIST_HEURISTIC")); + HeuristicPriorities = parseHeuristic( + schedIni.GetString("LIST_HEURISTIC", schedIni.GetString("HEURISTIC", "LLVM"))); EnumPriorities = parseHeuristic(schedIni.GetString("ENUM_HEURISTIC")); SecondPassEnumPriorities = parseHeuristic(schedIni.GetString("SECOND_PASS_ENUM_HEURISTIC")); @@ -1014,12 +1099,13 @@ void ScheduleDAGOptSched::loadOptSchedConfig() { RandomGen::SetSeed(randomSeed); HeurSchedType = parseListSchedType(); - OccupancyLimit = schedIni.GetInt("OCCUPANCY_LIMIT"); - ShouldLimitOccupancy = schedIni.GetBool("SHOULD_LIMIT_OCCUPANCY"); + OccupancyLimit = schedIni.GetInt("OCCUPANCY_LIMIT", 0); + ShouldLimitOccupancy = schedIni.GetBool("SHOULD_LIMIT_OCCUPANCY", false); OccupancyLimitSource = OCC_LIMIT_TYPE::OLT_NONE; if (ShouldLimitOccupancy) - OccupancyLimitSource = parseOccLimit(schedIni.GetString("OCCUPANCY_LIMIT_SOURCE")); + OccupancyLimitSource = + parseOccLimit(schedIni.GetString("OCCUPANCY_LIMIT_SOURCE", "NONE")); DeviceACOEnabled = schedIni.GetBool("DEV_ACO"); } From 88b7abbe931b0fa3c800dde0f6458cbb3b566e34 Mon Sep 17 00:00:00 2001 From: Peng Kang Date: Thu, 28 May 2026 19:12:26 +0000 Subject: [PATCH 29/30] Update - remove the unnecessary update and change print to Logger:Info MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Bug 1 fix — numAntsTerminated_ data race: FindManyCPUSchedule (aco.hip.cpp:2469): introduces a std::atomic localAntsTerminated{0}, captures it by reference in each thread lambda, and adds it to numAntsTerminated_ after joining. PCPU_FindOneSchedule signature (aco.h:157, aco.hip.cpp:2517): takes std::atomic &antsTerminated instead of touching the shared numAntsTerminated_ directly. The early-exit increment (aco.hip.cpp:2638): changed from numAntsTerminated_++ to antsTerminated.fetch_add(1, std::memory_order_relaxed). Bug 2 fix — shared RNG serializes all ants: PCPUACOSchedVars (aco.h:94): added std::mt19937 rng member. AllocPCPUACOSchedVars (aco.hip.cpp:2921): seeds each thread's RNG with random_seed_ ^ (thread_index * large_prime) so each ant gets a distinct, reproducible sequence. PCPU_SelectInstruction (aco.hip.cpp:2804): replaced RandDouble(...) calls (which contend on a global mutex) with std::uniform_real_distribution draws from pcpu_sched_vars.rng — fully lock-free per thread. --- example/optsched-cfg/sched.ini | 4 +- include/opt-sched/Scheduler/aco.h | 3 ++ lib/Scheduler/aco.hip.cpp | 55 ++++++++++--------- lib/Scheduler/sched_region.hip.cpp | 21 ++++---- lib/Wrapper/OptimizingScheduler.hip.cpp | 72 +++++-------------------- 5 files changed, 60 insertions(+), 95 deletions(-) diff --git a/example/optsched-cfg/sched.ini b/example/optsched-cfg/sched.ini index b472746d..dda3bb88 100644 --- a/example/optsched-cfg/sched.ini +++ b/example/optsched-cfg/sched.ini @@ -25,7 +25,7 @@ USE_TWO_PASS NO HEUR_ENABLED YES # ACO_ENABLED is the Ant Colony Optimization scheduler. ACO_ENABLED YES -o# MO_ENABLED toggles multiple occupancy schedule generation. +# MO_ENABLED toggles multiple occupancy schedule generation. # If YES, scheduler explores multiple occupancy points; otherwise it uses a single occupancy. MO_ENABLED NO # ENUM_ENABLED is the Branch and Bound scheduler. @@ -344,3 +344,5 @@ HIST_TABLE_HASH_BITS 16 CYCLES_FROM_OPTIMAL 0 # Threshold of when to skip ACO, determined by how deeply inserted in a loop list schedule is (-1 to disable this filter) LOOP_DEPTH -1 + +RANDOM_SEED 0 diff --git a/include/opt-sched/Scheduler/aco.h b/include/opt-sched/Scheduler/aco.h index cf21516a..0ab40580 100644 --- a/include/opt-sched/Scheduler/aco.h +++ b/include/opt-sched/Scheduler/aco.h @@ -15,6 +15,7 @@ Last Update: Jan. 2020 #include "opt-sched/Scheduler/device_vector.h" #include "llvm/ADT/ArrayRef.h" #include +#include #include #include #include @@ -91,6 +92,7 @@ struct alignas(64) PCPUACOSchedVars { pheromone_t MaxPriorityInv; KeysHelper1 *kHelper1; + std::mt19937 rng; /* void ConstrainedScheduler::SchdulInst_(SchedInstruction *inst, InstCount) @@ -156,6 +158,7 @@ class ACOScheduler : public ConstrainedScheduler { int thread, PCPUACOSchedVars &pcpu_sched_vars, ParallelCPUVars &pcpu, + std::atomic &antsTerminated, int kernelNum = -1); InstCount PCPU_SelectInstruction(SchedInstruction *lastInst, InstCount totalStalls, SchedRegion *rgn, bool &unnecessarilyStalling, diff --git a/lib/Scheduler/aco.hip.cpp b/lib/Scheduler/aco.hip.cpp index 5d14113f..c9f285eb 100644 --- a/lib/Scheduler/aco.hip.cpp +++ b/lib/Scheduler/aco.hip.cpp @@ -17,6 +17,8 @@ #include #include +#include +#include #include #include @@ -85,7 +87,7 @@ ACOScheduler::ACOScheduler(DataDepGraph *dataDepGraph, use_dev_ACO = schedIni.GetBool("DEV_ACO") && dev_rgn_ != nullptr && dev_DDG_ != nullptr; if (!use_dev_ACO || count_ < REGION_MIN_SIZE) { - numThreads_ = schedIni.GetInt("HOST_ANTS", 2); + numThreads_ = schedIni.GetInt("HOST_ANTS"); } else { dev_rgn_->SetNumThreads(numThreads_); dev_DDG_->SetNumThreads(numThreads_); @@ -97,7 +99,7 @@ ACOScheduler::ACOScheduler(DataDepGraph *dataDepGraph, local_decay = schedIni.GetFloat("ACO_LOCAL_DECAY"); print_aco_trace = schedIni.GetBool("ACO_TRACE"); IsTwoPassEn = schedIni.GetBool("USE_TWO_PASS"); - weightedSecondPass = schedIni.GetBool("USE_WEIGHTED_SECOND_PASS", false); + weightedSecondPass = schedIni.GetBool("USE_WEIGHTED_SECOND_PASS"); /* std::cerr << "useOldAlg===="<noImprovementMax = noImprovementMax; @@ -1427,13 +1429,13 @@ FUNC_RESULT ACOScheduler::FindSchedule(InstSchedule *schedule_out, #else initialValue_ = (double)numThreads_ / heuristicCost; #endif - #ifdef MULTIPLE_PHEROMONE_TABLES +#ifdef MULTIPLE_PHEROMONE_TABLES for (int i = 0; i < pheromone_size * numDiffOccupancies_; i++) pheromone_[i] = initialValue_; - #else +#else for (int i = 0; i < pheromone_size; i++) pheromone_[i] = initialValue_; - #endif +#endif InstSchedule *bestSchedule = InitialSchedule; // check if heuristic schedule is better than the initial @@ -2020,12 +2022,10 @@ void ACOScheduler::UpdatePheromone(InstSchedule *schedule, bool isIterationBest, #if !USE_ACS // decay pheromone (only if count_ is sane) - if (count_ > 0) { - for (int i = 0; i < count_; i++) { - for (int j = 0; j < count_; j++) { - pheromone = &Pheromone(i, j); - *pheromone *= (1 - decay_factor); - } + for (int i = 0; i < count_; i++) { + for (int j = 0; j < count_; j++) { + pheromone = &Pheromone(i, j); + *pheromone *= (1 - decay_factor); } } #endif @@ -2467,18 +2467,21 @@ InstSchedule *ACOScheduler::FindManyCPUSchedule(InstCount RPTarget) { std::vector PCPUThreads; InstSchedule **cpuScheds = new InstSchedule*[NO_CPU_THREADS](); ((BBWithSpill*)rgn_)->AllocParallelCPUVars(NO_CPU_THREADS); + std::atomic localAntsTerminated{0}; for (int i = 0; i < NO_CPU_THREADS; i++) { - PCPUThreads.emplace_back([this, cpuScheds, i, RPTarget, pcpu_sched_vars]() { - cpuScheds[i] = PCPU_FindOneSchedule(RPTarget, i, + PCPUThreads.emplace_back([this, cpuScheds, i, RPTarget, pcpu_sched_vars, &localAntsTerminated]() { + cpuScheds[i] = PCPU_FindOneSchedule(RPTarget, i, pcpu_sched_vars[i], - ((BBWithSpill*)rgn_)->GetPCPUVars(i)); + ((BBWithSpill*)rgn_)->GetPCPUVars(i), + localAntsTerminated); }); } for (auto &t : PCPUThreads) { t.join(); } + numAntsTerminated_ += localAntsTerminated.load(); ((BBWithSpill*)rgn_)->FreeParallelCPUVars(NO_CPU_THREADS); // Choose the best schedule from the CPU threads. @@ -2509,11 +2512,12 @@ InstSchedule *ACOScheduler::FindManyCPUSchedule(InstCount RPTarget) { return result; } -InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, +InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, int thread, PCPUACOSchedVars &pcpu_sched_vars, ParallelCPUVars &pcpu, - int kernelNum){ + std::atomic &antsTerminated, + int kernelNum){ Logger::Info("FindOneSchedule, %d", thread); SchedInstruction *lastInst = NULL; ACOReadyListEntry LastInstInfo; @@ -2635,7 +2639,7 @@ InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, // schedule construction if (((BBWithSpill*)rgn_)->PCPU_GetCrntSpillCost(pcpu) > RPTarget) { // end schedule construction - numAntsTerminated_++; + antsTerminated.fetch_add(1, std::memory_order_relaxed); pcpu_sched_vars.readyLs->clearReadyList(); delete schedule; Logger::Info("Thread %d, returned early", thread); @@ -2801,10 +2805,10 @@ InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstC //we are going to use the fixed bias or if we are going to use //fitness proportional selection. Generate the number used for //the fitness proportional selection point - double rand ; - pheromone_t point; - rand = RandDouble(0, 1); - point = RandDouble(0, pcpu_sched_vars.readyLs->ScoreSum); + std::uniform_real_distribution dist01(0.0, 1.0); + double rand = dist01(pcpu_sched_vars.rng); + pheromone_t point = std::uniform_real_distribution( + 0.0, pcpu_sched_vars.readyLs->ScoreSum)(pcpu_sched_vars.rng); //here we compute the chance that we will use fp selection or auto pick the best double choose_best_chance; @@ -2917,6 +2921,7 @@ PCPUACOSchedVars *ACOScheduler::AllocPCPUACOSchedVars(int numThreads) { pcpu_sched_vars[i].MaxPriorityInv = 0; pcpu_sched_vars[i].kHelper1 = new KeysHelper1(priorities1_); pcpu_sched_vars[i].kHelper1->initForRegion(dataDepGraph_); + pcpu_sched_vars[i].rng.seed(random_seed_ ^ (uint64_t(i + 1) * 6364136223846793005ULL)); pcpu_sched_vars[i].rsrvSlots = new ReserveSlot[issuRate_]; pcpu_sched_vars[i].avlblSlotsInCrntCycle = new int16_t[issuTypeCnt_]; diff --git a/lib/Scheduler/sched_region.hip.cpp b/lib/Scheduler/sched_region.hip.cpp index 65f00050..805f5abe 100644 --- a/lib/Scheduler/sched_region.hip.cpp +++ b/lib/Scheduler/sched_region.hip.cpp @@ -254,25 +254,24 @@ FUNC_RESULT SchedRegion::FindOptimalSchedule( Config &schedIni = SchedulerOptions::getInstance(); bool HeuristicSchedulerEnabled = schedIni.GetBool("HEUR_ENABLED"); bool AcoSchedulerEnabled = schedIni.GetBool("ACO_ENABLED"); - multipleOccupancies = schedIni.GetBool("MO_ENABLED", false); - printf("multipleOccupancies is %d", multipleOccupancies); + multipleOccupancies = schedIni.GetBool("MO_ENABLED"); bool BbSchedulerEnabled = isBbEnabled(schedIni, rgnTimeout); - unsigned long randSeed = (unsigned long) schedIni.GetInt("RANDOM_SEED", 0); - bool devACOEnabled = schedIni.GetBool("DEV_ACO", false); + unsigned long randSeed = (unsigned long) schedIni.GetInt("RANDOM_SEED"); + bool devACOEnabled = schedIni.GetBool("DEV_ACO"); int numBlocks; if (devACOEnabled && dataDepGraph_->GetInstCnt() >= REGION_MIN_SIZE) { - bool manyAntsEnabled = schedIni.GetBool("ACO_MANY_ANTS_ENABLED", false); - int manyAntBlocks = static_cast(schedIni.GetInt("ACO_MANY_ANTS_PER_ITERATION_BLOCKS", 1)); - int deviceAntBlocks = static_cast(schedIni.GetInt("ACO_DEVICE_ANT_PER_ITERATION_BLOCKS", 1)); + bool manyAntsEnabled = schedIni.GetBool("ACO_MANY_ANTS_ENABLED"); + int manyAntBlocks = static_cast(schedIni.GetInt("ACO_MANY_ANTS_PER_ITERATION_BLOCKS")); + int deviceAntBlocks = static_cast(schedIni.GetInt("ACO_DEVICE_ANT_PER_ITERATION_BLOCKS")); numBlocks = manyAntsEnabled && dataDepGraph_->GetInstCnt() > MANY_ANT_MIN_SIZE ? manyAntBlocks : deviceAntBlocks; } else { - numBlocks = static_cast(schedIni.GetInt("HOST_ANTS", 2)); + numBlocks = static_cast(schedIni.GetInt("HOST_ANTS")); } if (AcoSchedulerEnabled) { - AcoBeforeEnum = schedIni.GetBool("ACO_BEFORE_ENUM", false); - AcoAfterEnum = schedIni.GetBool("ACO_AFTER_ENUM", false); + AcoBeforeEnum = schedIni.GetBool("ACO_BEFORE_ENUM"); + AcoAfterEnum = schedIni.GetBool("ACO_AFTER_ENUM"); } if (!HeuristicSchedulerEnabled && !AcoBeforeEnum) { @@ -392,7 +391,7 @@ FUNC_RESULT SchedRegion::FindOptimalSchedule( CmputNormCost_(lstSched, CCM_DYNMC, hurstcExecCost, true); hurstcCost_ = lstSched->GetCost(); InstCount maxIndependentInstructions = 0; - std::string readyListUB = schedIni.GetString("ACO_READY_LIST_UB", "NO"); + std::string readyListUB = schedIni.GetString("ACO_READY_LIST_UB"); if (readyListUB == "NO" || readyListUB == "MIN_DEGREE") { for (int i = 0; i < dataDepGraph_->GetInstCnt(); i++) { int independentInstructions = dataDepGraph_->GetInstCnt() - dataDepGraph_->GetInstByIndx(i)->GetRcrsvPrdcsrCnt() - dataDepGraph_->GetInstByIndx(i)->GetRcrsvScsrCnt(); diff --git a/lib/Wrapper/OptimizingScheduler.hip.cpp b/lib/Wrapper/OptimizingScheduler.hip.cpp index acd159b9..5b2e4f9f 100644 --- a/lib/Wrapper/OptimizingScheduler.hip.cpp +++ b/lib/Wrapper/OptimizingScheduler.hip.cpp @@ -290,32 +290,14 @@ void ScheduleEvaluator::recordSchedule(int schedIndex) { if (!DAG.BB || DAG.RegionBegin == DAG.RegionEnd) return; - dbgs() << "[OPTSCHED_DEBUG] recordSchedule idx=" << schedIndex - << " this=" << (const void *)this - << " DAG=" << (const void *)&DAG - << " DAG.BB=" << (const void *)DAG.BB - << " RegionBegin==RegionEnd=" << (DAG.RegionBegin == DAG.RegionEnd) - << " storedSchedules.size()=" << storedSchedules.size() << "\n"; - if ((size_t)schedIndex >= storedSchedules.size()) storedSchedules.resize((size_t)schedIndex + 1); - dbgs() << "[OPTSCHED_DEBUG] recordSchedule idx=" << schedIndex - << " DAG.BB=" << (void *)DAG.BB - << " RegionBegin==RegionEnd=" << (DAG.RegionBegin == DAG.RegionEnd) - << " storedSchedules.size()=" << storedSchedules.size() << "\n"; - storedSchedules[schedIndex].clear(); schedCount++; const char *dbgEnv = std::getenv("OPTSCHED_DEBUG"); const bool doDebug = dbgEnv && std::atoi(dbgEnv) != 0; - if (doDebug) { - dbgs() << "[OPTSCHED_DEBUG] recordSchedule idx=" << schedIndex - << " DAG.BB=" << (void *)DAG.BB - << " RegionBegin==RegionEnd=" << (DAG.RegionBegin == DAG.RegionEnd) - << " storedSchedules.size()=" << storedSchedules.size() << "\n"; - } // Safely iterate over scheduled instructions in the region if (!DAG.BB || DAG.RegionBegin == DAG.RegionEnd) @@ -326,17 +308,6 @@ void ScheduleEvaluator::recordSchedule(int schedIndex) { storedSchedules[schedIndex].push_back(&*I); } } - - if (doDebug) { - dbgs() << "[OPTSCHED_DEBUG] storedSchedules[" << schedIndex << "] size=" - << storedSchedules[schedIndex].size() << "\n"; - int __mi_i = 0; - for (MachineInstr *MI : storedSchedules[schedIndex]) { - dbgs() << " [" << __mi_i++ << "] "; - MI->print(dbgs()); - dbgs() << '\n'; - } - } } // revert to a previous schedule @@ -744,12 +715,6 @@ void ScheduleDAGOptSched::schedule() { // Revord MachineInstr order in the first pass for a possible revert if // scheduling makes things worse. if (!SecondPass) { - LLVM_DEBUG(dbgs() << "[OPTSCHED_DEBUG] before first recordSchedule this=" - << (const void *)this - << " RegionNumber=" << RegionNumber - << " BB=" << (const void *)BB - << " RegionBegin==RegionEnd=" << (RegionBegin == RegionEnd) - << "\n"); SchedEval.recordSchedule(0); SchedEval.calculateRPBefore(); SchedEval.calcualteILPBefore(); @@ -783,25 +748,16 @@ void ScheduleDAGOptSched::schedule() { // Prepare for device scheduling by increasing heap size and copying machMdl bool dev_ACOEnabled = schedIni.GetBool("DEV_ACO"); if (dev_ACOEnabled && dev_MM == NULL && NumRegionInstrs + 2 >= REGION_MIN_SIZE) { - int hipDeviceCount = 0; - hipError_t hipErr = hipGetDeviceCount(&hipDeviceCount); - if (hipErr != hipSuccess || hipDeviceCount == 0) { - LLVM_DEBUG(dbgs() << "[OPTSCHED_DEBUG] DEV_ACO disabled: no HIP devices found (" - << hipGetErrorString(hipErr) << ")\n"); - dev_ACOEnabled = false; - } else { - gpuErrchk(hipSetDevice(0)); - // Copy MachineModel to device for use during DevListSched. - // Allocate device memory - gpuErrchk(hipMallocManaged((void**)&dev_MM, sizeof(MachineModel))); - // Copy machMdl_ to device - gpuErrchk(hipMemcpy(dev_MM, MM.get(), sizeof(MachineModel), - hipMemcpyHostToDevice)); - // Copy over all pointers to device - MM.get()->CopyPointersToDevice(dev_MM); - // make sure mallocmanaged mem is copied to device before kernel start - gpuErrchk(hipMemPrefetchAsync(dev_MM, sizeof(MachineModel), 0)); - } + // Copy MachineModel to device for use during DevListSched. + // Allocate device memory + gpuErrchk(hipMallocManaged((void**)&dev_MM, sizeof(MachineModel))); + // Copy machMdl_ to device + gpuErrchk(hipMemcpy(dev_MM, MM.get(), sizeof(MachineModel), + hipMemcpyHostToDevice)); + // Copy over all pointers to device + MM.get()->CopyPointersToDevice(dev_MM); + // make sure mallocmanaged mem is copied to device before kernel start + gpuErrchk(hipMemPrefetchAsync(dev_MM, sizeof(MachineModel), 0)); } // create region @@ -1073,7 +1029,7 @@ void ScheduleDAGOptSched::loadOptSchedConfig() { SCW = schedIni.GetInt("SPILL_COST_WEIGHT"); LowerBoundAlgorithm = parseLowerBoundAlgorithm(); HeuristicPriorities = parseHeuristic( - schedIni.GetString("LIST_HEURISTIC", schedIni.GetString("HEURISTIC", "LLVM"))); + schedIni.GetString("LIST_HEURISTIC")); EnumPriorities = parseHeuristic(schedIni.GetString("ENUM_HEURISTIC")); SecondPassEnumPriorities = parseHeuristic(schedIni.GetString("SECOND_PASS_ENUM_HEURISTIC")); @@ -1099,13 +1055,13 @@ void ScheduleDAGOptSched::loadOptSchedConfig() { RandomGen::SetSeed(randomSeed); HeurSchedType = parseListSchedType(); - OccupancyLimit = schedIni.GetInt("OCCUPANCY_LIMIT", 0); - ShouldLimitOccupancy = schedIni.GetBool("SHOULD_LIMIT_OCCUPANCY", false); + OccupancyLimit = schedIni.GetInt("OCCUPANCY_LIMIT"); + ShouldLimitOccupancy = schedIni.GetBool("SHOULD_LIMIT_OCCUPANCY"); OccupancyLimitSource = OCC_LIMIT_TYPE::OLT_NONE; if (ShouldLimitOccupancy) OccupancyLimitSource = - parseOccLimit(schedIni.GetString("OCCUPANCY_LIMIT_SOURCE", "NONE")); + parseOccLimit(schedIni.GetString("OCCUPANCY_LIMIT_SOURCE")); DeviceACOEnabled = schedIni.GetBool("DEV_ACO"); } From aee700e5bb2e4513efd92cbe8adaaf84983a41ac Mon Sep 17 00:00:00 2001 From: Peng Kang Date: Fri, 29 May 2026 03:15:27 +0000 Subject: [PATCH 30/30] feat: replace per-iteration thread spawn with persistent CPU thread pool Replace FindManyCPUSchedule's per-batch thread create/join with a persistent CPUThreadPool that lives for the duration of each occupancy iteration, eliminating the spawn/join overhead on every ACO iteration. Threads are kept alive between iterations using a generation-counter condition variable pattern. All per-thread state (SchedInstruction pcpu_inst_vars_, Register pcpu_crntUseCnt_, ParallelCPUVars fields and bitvectors, and the shared Register::crntUseCnt_ that IsLive() reads) is explicitly reset at the top of each iteration via new Reset methods, replacing the implicit reset that Free+Alloc in FindManyCPUSchedule previously provided. New methods added: - CPUThreadPool: persistent pool struct (aco.hip.cpp) - ACOScheduler::PCPU_ResetSchedInsts: resets pcpu_inst_vars_ per iter - SchedInstruction::InitPCPUVars: implements previously empty reset body - Register::ResetParallelCPURegs: zeroes per-thread crntUseCnt_ array - BBWithSpill::ResetParallelCPUVars: resets all ParallelCPUVars state including shared crntUseCnt_ to fix "used without being defined" fatal --- include/opt-sched/Scheduler/aco.h | 9 + include/opt-sched/Scheduler/bb_spill.h | 1 + include/opt-sched/Scheduler/register.h | 1 + lib/Scheduler/aco.hip.cpp | 222 ++++++++++++++++++++----- lib/Scheduler/bb_spill.hip.cpp | 32 ++++ lib/Scheduler/register.hip.cpp | 5 + lib/Scheduler/sched_basic_data.hip.cpp | 21 +++ 7 files changed, 251 insertions(+), 40 deletions(-) diff --git a/include/opt-sched/Scheduler/aco.h b/include/opt-sched/Scheduler/aco.h index 0ab40580..ba2deedc 100644 --- a/include/opt-sched/Scheduler/aco.h +++ b/include/opt-sched/Scheduler/aco.h @@ -14,12 +14,19 @@ Last Update: Jan. 2020 #include "opt-sched/Scheduler/ready_list.h" #include "opt-sched/Scheduler/device_vector.h" #include "llvm/ADT/ArrayRef.h" +#include +#include +#include #include +#include #include +#include #include #include #include +struct CPUThreadPool; + namespace llvm { namespace opt_sched { @@ -182,6 +189,7 @@ class ACOScheduler : public ConstrainedScheduler { bool PCPU_MovToNxtSlot_(SchedInstruction *inst, PCPUACOSchedVars &pcpu_sched_vars); void PCPU_InitNewCycle_(PCPUACOSchedVars &pcpu_sched_vars); void PCPU_InitSchedInsts(int numThreads); + void PCPU_ResetSchedInsts(int numThreads); void PCPU_FreeSchedInsts(int numThreads); __host__ __device__ @@ -304,6 +312,7 @@ class ACOScheduler : public ConstrainedScheduler { // Populated before PCPU threads launch to avoid cursor race on scsrLst_. std::vector>> instScsrs_; + std::unique_ptr cpuPool_; }; } // namespace opt_sched diff --git a/include/opt-sched/Scheduler/bb_spill.h b/include/opt-sched/Scheduler/bb_spill.h index 323f91f3..99d0ae31 100644 --- a/include/opt-sched/Scheduler/bb_spill.h +++ b/include/opt-sched/Scheduler/bb_spill.h @@ -272,6 +272,7 @@ class BBWithSpill : public SchedRegion { // } void AllocParallelCPUVars(int numThreads); + void ResetParallelCPUVars(int numThreads); void FreeParallelCPUVars(int numThreads); void PCPU_CmputCrntSpillCost_(ParallelCPUVars &pcpu); diff --git a/include/opt-sched/Scheduler/register.h b/include/opt-sched/Scheduler/register.h index 0d200f81..451c6496 100644 --- a/include/opt-sched/Scheduler/register.h +++ b/include/opt-sched/Scheduler/register.h @@ -136,6 +136,7 @@ class Register { void PCPU_ResetCrntUseCnt(int threadIdx); bool PCPU_IsLive(int threadIdx) const; void AllocParallelCPURegs(int numThreads); + void ResetParallelCPURegs(int numThreads); void FreeParallelCPURegs(); diff --git a/lib/Scheduler/aco.hip.cpp b/lib/Scheduler/aco.hip.cpp index c9f285eb..272dcce6 100644 --- a/lib/Scheduler/aco.hip.cpp +++ b/lib/Scheduler/aco.hip.cpp @@ -25,6 +25,73 @@ using namespace llvm::opt_sched; namespace cg = cooperative_groups; +// Persistent thread pool: threads are created once and reused across ACO +// iterations, eliminating per-iteration spawn/join overhead. +struct CPUThreadPool { + void start(int n) { + shutdown_ = false; + generation_ = 0; + pending_ = 0; + for (int i = 0; i < n; i++) + workers_.emplace_back([this, i] { workerLoop(i); }); + } + + // Fan out fn(threadIdx) to all workers and block until all finish. + void runAll(std::function fn) { + { + std::unique_lock lock(mtx_); + task_ = std::move(fn); + pending_ = static_cast(workers_.size()); + ++generation_; + } + startCv_.notify_all(); + std::unique_lock lock(mtx_); + doneCv_.wait(lock, [this] { return pending_ == 0; }); + } + + void stop() { + { + std::unique_lock lock(mtx_); + shutdown_ = true; + ++generation_; + } + startCv_.notify_all(); + for (auto &t : workers_) + t.join(); + workers_.clear(); + } + +private: + void workerLoop(int idx) { + int seenGen = 0; + for (;;) { + std::function fn; + { + std::unique_lock lock(mtx_); + startCv_.wait(lock, + [this, &seenGen] { return generation_ != seenGen || shutdown_; }); + if (shutdown_) + return; + seenGen = generation_; + fn = task_; + } + fn(idx); + std::unique_lock lock(mtx_); + if (--pending_ == 0) + doneCv_.notify_one(); + } + } + + std::vector workers_; + std::function task_; + std::mutex mtx_; + std::condition_variable startCv_; + std::condition_variable doneCv_; + int generation_{0}; + int pending_{0}; + bool shutdown_{false}; +}; + #ifndef NDEBUG static void PrintInstruction(SchedInstruction *inst); #endif @@ -1802,53 +1869,112 @@ FUNC_RESULT ACOScheduler::FindSchedule(InstSchedule *schedule_out, std::unordered_map schedMap; int diffSchedCount = 0; #endif - - while (noImprovement < noImprovementMax) { + + // One-time setup per occupancy iteration: allocate per-thread state and + // build the successor cache so threads never race on cursor state. + PCPU_InitSchedInsts(numThreads_); + PCPUACOSchedVars *pcpu_sched_vars = AllocPCPUACOSchedVars(numThreads_); + instScsrs_.assign(count_, {}); + for (int si = 0; si < count_; si++) { + SchedInstruction *sInst = dataDepGraph_->GetInstByIndx(si); + if (!sInst) continue; + InstCount prdcsrNum; + for (SchedInstruction *scsr = sInst->GetFrstScsr(&prdcsrNum); + scsr != NULL; scsr = sInst->GetNxtScsr(&prdcsrNum)) + instScsrs_[si].emplace_back(scsr, prdcsrNum); + } + InstSchedule **cpuScheds = new InstSchedule *[numThreads_](); + ((BBWithSpill *)rgn_)->AllocParallelCPUVars(numThreads_); + std::atomic localAntsTerminated{0}; + cpuPool_ = std::make_unique(); + cpuPool_->start(numThreads_); + + while (noImprovement < noImprovementMax) { iterations++; iterationBest = nullptr; - for (int i = 0; i < numThreads_; i++) { - InstSchedule *schedule; - if (!RUN_PCPU) { - schedule = FindOneSchedule(RPTarget, NULL); - } else { - schedule = FindManyCPUSchedule(RPTarget); - Logger::Info("Done Find Many Schedule"); - i+=(NO_CPU_THREADS-1); + + if (!RUN_PCPU) { + // Sequential fallback path. + for (int i = 0; i < numThreads_; i++) { + InstSchedule *schedule = FindOneSchedule(RPTarget, NULL); + if (print_aco_trace) + PrintSchedule(schedule); + if (shouldReplaceSchedule(iterationBest, schedule, false, RPTarget, targetOccupancy_ - j)) { + if (iterationBest) + delete iterationBest; + iterationBest = schedule; + } else { + if (schedule) + delete schedule; + } + } + } else { + // Parallel path: dispatch all ants at once via persistent pool. + // Reset all per-iteration state that was previously zeroed by the + // Free+Alloc cycle inside FindManyCPUSchedule. + PCPU_ResetSchedInsts(numThreads_); + ((BBWithSpill *)rgn_)->ResetParallelCPUVars(numThreads_); + for (int i = 0; i < numThreads_; i++) { + pcpu_sched_vars[i].schduldInstCnt = 0; + pcpu_sched_vars[i].isCrntCycleBlkd = false; + pcpu_sched_vars[i].crntCycleNum = 0; + pcpu_sched_vars[i].crntSlotNum = 0; + pcpu_sched_vars[i].rsrvSlotCnt = 0; + pcpu_sched_vars[i].RP0OrPositiveCount = 0; + pcpu_sched_vars[i].MaxScoringInst = 0; + pcpu_sched_vars[i].readyLs->clearReadyList(); + for (int j = 0; j < issuRate_; j++) { + pcpu_sched_vars[i].rsrvSlots[j].strtCycle = INVALID_VALUE; + pcpu_sched_vars[i].rsrvSlots[j].endCycle = INVALID_VALUE; + } + for (int j = 0; j < issuTypeCnt_; j++) + pcpu_sched_vars[i].avlblSlotsInCrntCycle[j] = slotsPerTypePerCycle_[j]; } + std::fill(cpuScheds, cpuScheds + numThreads_, nullptr); + cpuPool_->runAll([&](int i) { + cpuScheds[i] = PCPU_FindOneSchedule(RPTarget, i, + pcpu_sched_vars[i], + ((BBWithSpill *)rgn_)->GetPCPUVars(i), + localAntsTerminated); + }); + + for (int i = 0; i < numThreads_; i++) { + InstSchedule *schedule = cpuScheds[i]; + if (!schedule) + continue; + cpuScheds[i] = nullptr; - #ifdef CHECK_DIFFERENT_SCHEDULES - // check if schedule is in Map - InstCount instNum, cycleNum, slotNum; - // get first instruction in string - instNum = schedule->GetFrstInst(cycleNum, slotNum); - std::string schedString = std::to_string(instNum); - // prepare next instruction in comma separated list - instNum = schedule->GetNxtInst(cycleNum, slotNum); - while (instNum != INVALID_VALUE) { - schedString.append(","); - schedString.append(std::to_string(instNum)); + #ifdef CHECK_DIFFERENT_SCHEDULES + InstCount instNum, cycleNum, slotNum; + instNum = schedule->GetFrstInst(cycleNum, slotNum); + std::string schedString = std::to_string(instNum); instNum = schedule->GetNxtInst(cycleNum, slotNum); + while (instNum != INVALID_VALUE) { + schedString.append(","); + schedString.append(std::to_string(instNum)); + instNum = schedule->GetNxtInst(cycleNum, slotNum); + } + schedule->ResetInstIter(); + if (schedMap.find(schedString) == schedMap.end()) { + schedMap[schedString] = 1; + diffSchedCount++; + } else { + schedMap[schedString] = schedMap[schedString] + 1; + } + #endif + + if (print_aco_trace) + PrintSchedule(schedule); + if (shouldReplaceSchedule(iterationBest, schedule, false, RPTarget, targetOccupancy_ - j)) { + if (iterationBest) + delete iterationBest; + iterationBest = schedule; + MaxPriorityInv = pcpu_sched_vars[i].MaxPriorityInv; + } else { + if (schedule) + delete schedule; } - schedule->ResetInstIter(); - if (schedMap.find(schedString) == schedMap.end()) { - schedMap[schedString] = 1; - diffSchedCount++; - } - else { - schedMap[schedString] = schedMap[schedString] + 1; - } - #endif - - if (print_aco_trace) - PrintSchedule(schedule); - if (shouldReplaceSchedule(iterationBest, schedule, false, RPTarget, targetOccupancy_ - j)) { - if (iterationBest) - delete iterationBest; - iterationBest = schedule; - } else { - if (schedule) - delete schedule; } } #if !USE_ACS @@ -1895,6 +2021,17 @@ FUNC_RESULT ACOScheduler::FindSchedule(InstSchedule *schedule_out, #endif } + + // Teardown per-occupancy-iteration parallel state. + cpuPool_->stop(); + cpuPool_.reset(); + numAntsTerminated_ += localAntsTerminated.load(); + ((BBWithSpill *)rgn_)->FreeParallelCPUVars(numThreads_); + FreePCPUACOSchedVars(pcpu_sched_vars, numThreads_); + pcpu_sched_vars = nullptr; + delete[] cpuScheds; + PCPU_FreeSchedInsts(numThreads_); + Logger::Info("%d ants terminated early", numAntsTerminated_); #ifdef CHECK_DIFFERENT_SCHEDULES Logger::Info("%d different schedules for %d total ants", diffSchedCount, (iterations + 1) * numThreads_ - numAntsTerminated_); @@ -3085,6 +3222,11 @@ void ACOScheduler::PCPU_InitSchedInsts(int numThreads){ } } +void ACOScheduler::PCPU_ResetSchedInsts(int numThreads) { + for (int i = 0; i < totInstCnt_; i++) + dataDepGraph_->GetInstByIndx(i)->InitPCPUVars(numThreads); +} + void ACOScheduler::PCPU_FreeSchedInsts(int numThreads){ for (int i = 0; i < totInstCnt_; i++){ SchedInstruction *inst = dataDepGraph_->GetInstByIndx(i); diff --git a/lib/Scheduler/bb_spill.hip.cpp b/lib/Scheduler/bb_spill.hip.cpp index 2c51b9c6..cdcd18e6 100644 --- a/lib/Scheduler/bb_spill.hip.cpp +++ b/lib/Scheduler/bb_spill.hip.cpp @@ -2029,6 +2029,38 @@ void BBWithSpill::AllocParallelCPUVars(int numThreads){ } } +void BBWithSpill::ResetParallelCPUVars(int numThreads) { + for (int i = 0; i < numThreads; i++) { + pcpu_vars_[i].crntCycleNum_ = 0; + pcpu_vars_[i].crntSlotNum_ = 0; + pcpu_vars_[i].crntSpillCost_ = 0; + pcpu_vars_[i].crntStepNum_ = -1; + pcpu_vars_[i].peakSpillCost_ = 0; + pcpu_vars_[i].totSpillCost_ = 0; + pcpu_vars_[i].slilSpillCost_ = 0; + pcpu_vars_[i].dynamicSlilLowerBound_ = staticSlilLowerBound_; + pcpu_vars_[i].schduldInstCnt_ = 0; + pcpu_vars_[i].schduldEntryInstCnt_ = 0; + pcpu_vars_[i].schduldExitInstCnt_ = 0; + for (int j = 0; j < regTypeCnt_; j++) { + pcpu_vars_[i].liveRegs_[j].Reset(); + pcpu_vars_[i].livePhysRegs_[j].Reset(); + pcpu_vars_[i].peakRegPressures_[j] = 0; + pcpu_vars_[i].regPressures_[j] = 0; + pcpu_vars_[i].sumOfLiveIntervalLengths_[j] = 0; + } + std::fill(pcpu_vars_[i].spillCosts_, + pcpu_vars_[i].spillCosts_ + dataDepGraph_->GetInstCnt(), 0); + } + for (int i = 0; i < regTypeCnt_; i++) + for (int j = 0; j < regFiles_[i].GetRegCnt(); j++) + regFiles_[i].GetReg(j)->ResetParallelCPURegs(numThreads); + // Reset the shared crntUseCnt_ that IsLive() reads. This mirrors what + // Initialize_() -> InitForCostCmputtn_() did in the old per-batch code. + for (int i = 0; i < regTypeCnt_; i++) + regFiles_[i].ResetCrntUseCnts(); +} + void BBWithSpill::FreeParallelCPUVars(int numThreads){ for (int i = 0; i < numThreads; i++) { delete[] pcpu_vars_[i].liveRegs_; diff --git a/lib/Scheduler/register.hip.cpp b/lib/Scheduler/register.hip.cpp index 23441c33..7022301e 100644 --- a/lib/Scheduler/register.hip.cpp +++ b/lib/Scheduler/register.hip.cpp @@ -217,6 +217,11 @@ void Register::AllocParallelCPURegs(int numThreads){ pcpu_crntUseCnt_[i].value = 0; } } +void Register::ResetParallelCPURegs(int numThreads) { + for (int i = 0; i < numThreads; i++) + pcpu_crntUseCnt_[i].value = 0; +} + void Register::FreeParallelCPURegs(){ delete[] pcpu_crntUseCnt_; pcpu_crntUseCnt_ = nullptr; diff --git a/lib/Scheduler/sched_basic_data.hip.cpp b/lib/Scheduler/sched_basic_data.hip.cpp index aeaa698d..874dd449 100644 --- a/lib/Scheduler/sched_basic_data.hip.cpp +++ b/lib/Scheduler/sched_basic_data.hip.cpp @@ -1382,6 +1382,27 @@ void SchedInstruction::AllocPCPUVars(int numThreads){ } } +void SchedInstruction::InitPCPUVars(int numThreads) { + if (!pcpu_inst_vars_) + return; + for (int i = 0; i < numThreads; ++i) { + pcpu_inst_vars_[i].crntSchedCycle = SCHD_UNSCHDULD; + pcpu_inst_vars_[i].crntSchedSlot = SCHD_UNSCHDULD; + pcpu_inst_vars_[i].lastUseCnt = 0; + pcpu_inst_vars_[i].ready = false; + pcpu_inst_vars_[i].minRdyCycle = INVALID_VALUE; + pcpu_inst_vars_[i].unschduldPrdcsrCnt = prdcsrCnt_; + pcpu_inst_vars_[i].unschduldScsrCnt = scsrCnt_; + pcpu_inst_vars_[i].crntRlxdCycle = SCHD_UNSCHDULD; + for (InstCount j = 0; j < prdcsrCnt_; ++j) { + if (pcpu_inst_vars_[i].rdyCyclePerPrdcsr) + pcpu_inst_vars_[i].rdyCyclePerPrdcsr[j] = INVALID_VALUE; + if (pcpu_inst_vars_[i].prevMinRdyCyclePerPrdcsr) + pcpu_inst_vars_[i].prevMinRdyCyclePerPrdcsr[j] = INVALID_VALUE; + } + } +} + void SchedInstruction::FreePCPUVars(int numThreads) { if (!pcpu_inst_vars_) return;