diff --git a/example/optsched-cfg/sched.ini b/example/optsched-cfg/sched.ini index 7db698d9..dda3bb88 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 +# 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. @@ -297,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 3fd688dd..ba2deedc 100644 --- a/include/opt-sched/Scheduler/aco.h +++ b/include/opt-sched/Scheduler/aco.h @@ -14,9 +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 { @@ -50,6 +60,77 @@ 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_; + + 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* +}; + + +struct alignas(64) PCPUACOSchedVars { + //ACOScheduler + ACOReadyList *readyLs; + InstCount MaxScoringInst; + int RP0OrPositiveCount; + + //ConstrainedScheduler:InstSchedule + InstCount schduldInstCnt; + bool isCrntCycleBlkd; + InstCount crntCycleNum; + InstCount crntSlotNum; + int16_t rsrvSlotCnt; + ReserveSlot *rsrvSlots; + int16_t *avlblSlotsInCrntCycle; + + pheromone_t MaxPriorityInv; + KeysHelper1 *kHelper1; + std::mt19937 rng; + + /* + void ConstrainedScheduler::SchdulInst_(SchedInstruction *inst, InstCount) + schduldInstCnt_ + isCrntCycleBlkd_ + void ConstrainedScheduler::DoRsrvSlots_(SchedInstruction *inst) + rsrvSlots_ + rsrvSlotCnt_ + crntCycleNum_ + crntSlotNum_ + + void ConstrainedScheduler::UpdtSlotAvlblty_(SchedInstruction *inst) + avlblSlotsInCrntCycle_; + + isSchedComplete ??? -> requires investigation what is totInstCnt + + 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 +}; + class ACOScheduler : public ConstrainedScheduler { public: ACOScheduler(DataDepGraph *dataDepGraph, MachineModel *machineModel, @@ -77,6 +158,40 @@ 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); + InstSchedule *PCPU_FindOneSchedule(InstCount RPTarget, + 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, + 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, + 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); + 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_ResetSchedInsts(int numThreads); + void PCPU_FreeSchedInsts(int numThreads); + __host__ __device__ InstSchedule *FindOneSchedule(InstCount RPTarget, InstSchedule *dev_schedule = NULL, int kernelNum = -1); @@ -129,6 +244,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, @@ -190,6 +308,11 @@ 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_; + + 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 a2d807b1..99d0ae31 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,37 @@ class BBWithSpill : public SchedRegion { // size_t calculateMemoryNeeded() { // return regTypeCnt_ * sizeof(WeightedBitVector) * numThreads * 2; // } + + void AllocParallelCPUVars(int numThreads); + void ResetParallelCPUVars(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/graph.h b/include/opt-sched/Scheduler/graph.h index a1539020..8ca11726 100644 --- a/include/opt-sched/Scheduler/graph.h +++ b/include/opt-sched/Scheduler/graph.h @@ -628,13 +628,13 @@ inline GraphNode *GraphNode::GetNxtPrdcsr(UDT_GLABEL &label) { __host__ inline GraphNode *GraphNode::GetFrstScsr() { - UDT_GLABEL label; + UDT_GLABEL label{}; return GetFrstScsr(label); } __host__ inline GraphNode *GraphNode::GetNxtScsr() { - UDT_GLABEL label; + UDT_GLABEL label{}; return GetNxtScsr(label); } diff --git a/include/opt-sched/Scheduler/register.h b/include/opt-sched/Scheduler/register.h index 86ebf6d9..451c6496 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,14 @@ 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 ResetParallelCPURegs(int numThreads); + void FreeParallelCPURegs(); + + private: int16_t type_; int num_; @@ -136,6 +148,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/include/opt-sched/Scheduler/sched_basic_data.h b/include/opt-sched/Scheduler/sched_basic_data.h index 859f6d3d..e16177c9 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,20 @@ class SchedInstruction : public GraphNode { // parallel ACO on device void AllocDevArraysForParallelACO(int numThreads); + + 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; // This instruction's index in the scsrs_, latencies_, predOrder_ arrays @@ -741,6 +774,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 09aa9b83..272dcce6 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,17 +11,87 @@ #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 +#include +#include 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 @@ -51,6 +124,9 @@ double RandDouble(double min, double max) { //#define DECAY_FACTOR 0.5 //#endif +#define RUN_PCPU 1 +#define NO_CPU_THREADS 8 + ACOScheduler::ACOScheduler(DataDepGraph *dataDepGraph, MachineModel *machineModel, InstCount upperBound, SchedPriorities priorities1, SchedPriorities priorities2, bool vrfySched, @@ -76,10 +152,10 @@ ACOScheduler::ACOScheduler(DataDepGraph *dataDepGraph, numDiffOccupancies_ = numDiffOccupancies; targetOccupancy_ = targetOccupancy; - use_dev_ACO = schedIni.GetBool("DEV_ACO"); - if(!use_dev_ACO || count_ < REGION_MIN_SIZE) + 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"); - else { + } else { dev_rgn_->SetNumThreads(numThreads_); dev_DDG_->SetNumThreads(numThreads_); } @@ -299,7 +375,7 @@ InstCount ACOScheduler::SelectInstruction(SchedInstruction *lastInst, InstCount pheromone_t MaxScore = -1; InstCount MaxScoreIndx = 0; dev_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 @@ -412,7 +488,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 @@ -432,7 +508,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; @@ -450,14 +526,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 @@ -874,7 +952,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); @@ -1414,14 +1496,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 - std::cerr << "initialValue_" << initialValue_ << std::endl; +#endif InstSchedule *bestSchedule = InitialSchedule; // check if heuristic schedule is better than the initial @@ -1789,44 +1870,111 @@ FUNC_RESULT ACOScheduler::FindSchedule(InstSchedule *schedule_out, int diffSchedCount = 0; #endif + // 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 = FindOneSchedule(RPTarget, NULL); - - #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)); - instNum = schedule->GetNxtInst(cycleNum, slotNum); + + 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; } - schedule->ResetInstIter(); - if (schedMap.find(schedString) == schedMap.end()) { - schedMap[schedString] = 1; - diffSchedCount++; + } + } 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; } - else { - schedMap[schedString] = schedMap[schedString] + 1; + 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 + 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; } - #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 @@ -1871,7 +2019,19 @@ FUNC_RESULT ACOScheduler::FindSchedule(InstSchedule *schedule_out, #if USE_ACS UpdatePheromone(bestSchedule, false); #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_); @@ -1889,6 +2049,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) @@ -1959,6 +2120,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; @@ -1969,8 +2137,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 @@ -1988,7 +2158,7 @@ void ACOScheduler::UpdatePheromone(InstSchedule *schedule, bool isIterationBest, schedule->ResetInstIter(); #if !USE_ACS - // decay pheromone + // decay pheromone (only if count_ is sane) for (int i = 0; i < count_; i++) { for (int j = 0; j < count_; j++) { pheromone = &Pheromone(i, j); @@ -2171,14 +2341,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 @@ -2191,17 +2366,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 } @@ -2408,3 +2580,666 @@ void ACOScheduler::FreeDevicePointers(bool IsSecondPass) { hipFree(dev_kHelper2); hipFree(pheromone_.elmnts_); } + +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); + + // 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); + std::atomic localAntsTerminated{0}; + + for (int i = 0; i < NO_CPU_THREADS; 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), + localAntsTerminated); + }); + } + + for (auto &t : PCPUThreads) { + t.join(); + } + numAntsTerminated_ += localAntsTerminated.load(); + ((BBWithSpill*)rgn_)->FreeParallelCPUVars(NO_CPU_THREADS); + + // Choose the best schedule from the CPU threads. + InstSchedule *result = nullptr; + MaxPriorityInv = 0; + + 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); + + Logger::Info("FindManyCPUSchedule done"); + return result; +} + +InstSchedule *ACOScheduler::PCPU_FindOneSchedule(InstCount RPTarget, + int thread, + PCPUACOSchedVars &pcpu_sched_vars, + ParallelCPUVars &pcpu, + std::atomic &antsTerminated, + int kernelNum){ + Logger::Info("FindOneSchedule, %d", thread); + 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 = pcpu_sched_vars.kHelper1->getMaxValue(); + if (MaxPriority == 0) + MaxPriority = 1; // divide by 0 is bad + //Initialize_(); //For InstSchedule/ConstrainedSchedule + + SchedInstruction *waitFor = NULL; + InstCount waitUntil = 0; + 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 + Logger::Info("ReadyListFunctions 1, %d", thread); + InstCount RootId = rootInst_->GetNum(); + 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; + pcpu_sched_vars.MaxScoringInst = 0; + lastInst = dataDepGraph_->GetInstByIndx(RootId); + bool closeToRPTarget = false; + pcpu_sched_vars.RP0OrPositiveCount = 0; + + Logger::Info("While not PCPU_IsSchedComplete, %d", thread); + SchedInstruction *inst = NULL; + 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); + 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; + } + } + } + + // 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 + 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, + thread, pcpu_sched_vars, pcpu); + + if (SelIndx != -1) { + LastInstInfo = pcpu_sched_vars.readyLs->removeInstructionAtIndex(SelIndx); + + InstCount InstId = LastInstInfo.InstId; + inst = dataDepGraph_->GetInstByIndx(InstId); + // potentially wait on the current instruction + 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 + 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; + } + } + } + + //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 (PCPU_ChkInstLglty_(waitFor, pcpu_sched_vars)) { + inst = waitFor; + waitFor = NULL; + lastInst = inst; + } + } + + // boilerplate, mostly copied from ListScheduler, try not to touch it + //Logger::Info("BoilerPlate, %d", thread); + InstCount instNum; + if (!inst) { + instNum = SCHD_STALL; + schedule->incrementTotalStalls(); + if (unnecessarilyStalling) + schedule->incrementUnnecessaryStalls(); + } else { + instNum = inst->GetNum(); + PCPU_SchdulInst_(inst, pcpu_sched_vars); + //SchdulInst(inst, pcpu_sched_vars.crntCycleNum) + 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 + if (((BBWithSpill*)rgn_)->PCPU_GetCrntSpillCost(pcpu) > RPTarget) { + // end schedule construction + antsTerminated.fetch_add(1, std::memory_order_relaxed); + pcpu_sched_vars.readyLs->clearReadyList(); + delete schedule; + Logger::Info("Thread %d, returned early", thread); + return NULL; + } + PCPU_DoRsrvSlots_(inst, pcpu_sched_vars); + // this is annoying + 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 (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("Thread, %d, Done", thread); + return schedule; +} + +InstCount ACOScheduler::PCPU_SelectInstruction(SchedInstruction *lastInst, InstCount totalStalls, + SchedRegion *rgn, bool &unnecessarilyStalling, + 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 + if (currentlyWaiting && pcpu_sched_vars.RP0OrPositiveCount == 0) + return -1; + + // calculate MaxScoringInst, and ScoreSum + pheromone_t MaxScore = -1; + InstCount MaxScoreIndx = 0; + 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; + + for (InstCount I = 0; I < readyListSize; ++I) { + RPIsHigh = false; + // 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 = *heurPtr; + pheromone_t IScore = PCPU_Score(lastInstId, CandidateId, Heur, !rgn->IsSecondPass(), pcpu_sched_vars); + if (pcpu_sched_vars.RP0OrPositiveCount != 0 && candidateDefs > candidateLUC) + IScore *= 0.9; + + *scorePtr = IScore; + + if (currentlyWaiting) { + // 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 + 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 (*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 = *readyOnPtr - 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(*pcpu_sched_vars.readyLs->getInstIdAtIndex(I)); + if (!tempInst) + continue; + RegIndxTuple *uses; + Register *use; + uint16_t usesCount = tempInst->GetUses(uses); + 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; + } + } + } + + // 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; + *pcpu_sched_vars.readyLs->getInstScoreAtIndex(I) = IScore; + pcpu_sched_vars.readyLs->ScoreSum += IScore; + + 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 + //the fitness proportional selection point + 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; + 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 < 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 + 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; + else + unnecessarilyStalling = false; + + //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){ + 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. + // 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); + + 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 = 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 = 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 + //LUC component, and then compute the score + 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->PCPU_CmputLastUseCnt(dataDepGraph_->RegFiles, thread); + LUCVal <<= LUCEntry.Offset; + Heur &= LUCVal; + } + 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; + } + //Logger::Info("Finish UpdateACOReadyList, %d", thread); +} + +//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; + + //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].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_]; + + //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; + } + + for (int j = 0; j < issuTypeCnt_; j++) { + pcpu_sched_vars[i].avlblSlotsInCrntCycle[j] = slotsPerTypePerCycle_[j]; + } + } + return pcpu_sched_vars; +} + +void ACOScheduler::FreePCPUACOSchedVars(PCPUACOSchedVars *pcpu_sched_vars, + int numThreads) { + if (!pcpu_sched_vars) + return; + + for (int i = 0; i < numThreads; i++) { + 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; + + delete[] pcpu_sched_vars[i].avlblSlotsInCrntCycle; + pcpu_sched_vars[i].avlblSlotsInCrntCycle = nullptr; + } + + 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_(); //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; + + // 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)) { + 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++; + //Logger::Info("SchduldInstCnt Incremented, %d", 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) { + //Logger::Info("Checking if sched complete if %d == %d", pcpu_sched_vars.schduldInstCnt, totInstCnt_); + return pcpu_sched_vars.schduldInstCnt == totInstCnt_; +} + +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); +} + + +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; +} + +void ACOScheduler::PCPU_InitSchedInsts(int numThreads){ + for (int i = 0; i < totInstCnt_; i++){ + SchedInstruction *inst = dataDepGraph_->GetInstByIndx(i); + inst->AllocPCPUVars(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); + 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; +} diff --git a/lib/Scheduler/bb_spill.hip.cpp b/lib/Scheduler/bb_spill.hip.cpp index c3ed7890..cdcd18e6 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; @@ -557,6 +558,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 +896,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 +1171,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 +1366,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 +1979,175 @@ 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::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_; + 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]; +} 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 c998b460..2d6348b4 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 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 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 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 lock(m); GenerateNextNumber(); return randNum % (max - min + 1) + min; } uint32_t RandomGen::GetRand32() { + std::scoped_lock lock(m); GenerateNextNumber(); return randNum; } uint64_t RandomGen::GetRand64() { + std::scoped_lock 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 lock(m); uint16_t bytesNeeded = (bitCnt + 7) / 8; uint16_t index = 0; diff --git a/lib/Scheduler/register.hip.cpp b/lib/Scheduler/register.hip.cpp index c010ab29..7022301e 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,34 @@ 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::ResetParallelCPURegs(int numThreads) { + for (int i = 0; i < numThreads; i++) + pcpu_crntUseCnt_[i].value = 0; +} + +void Register::FreeParallelCPURegs(){ + delete[] pcpu_crntUseCnt_; + pcpu_crntUseCnt_ = nullptr; } __host__ 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_basic_data.hip.cpp b/lib/Scheduler/sched_basic_data.hip.cpp index 8bef6bdc..874dd449 100644 --- a/lib/Scheduler/sched_basic_data.hip.cpp +++ b/lib/Scheduler/sched_basic_data.hip.cpp @@ -1352,3 +1352,121 @@ __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::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; + + 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; +} + +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; +} + diff --git a/lib/Scheduler/sched_region.hip.cpp b/lib/Scheduler/sched_region.hip.cpp index 3a44f5d9..805f5abe 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; @@ -255,16 +255,19 @@ FUNC_RESULT SchedRegion::FindOptimalSchedule( bool HeuristicSchedulerEnabled = schedIni.GetBool("HEUR_ENABLED"); bool AcoSchedulerEnabled = schedIni.GetBool("ACO_ENABLED"); multipleOccupancies = schedIni.GetBool("MO_ENABLED"); - 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"); 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"); + 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")); + } if (AcoSchedulerEnabled) { AcoBeforeEnum = schedIni.GetBool("ACO_BEFORE_ENUM"); @@ -1143,6 +1146,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..5b2e4f9f 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,48 @@ 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; + + if ((size_t)schedIndex >= storedSchedules.size()) + storedSchedules.resize((size_t)schedIndex + 1); storedSchedules[schedIndex].clear(); - storedSchedules[schedIndex].reserve(DAG.NumRegionInstrs); + schedCount++; + + const char *dbgEnv = std::getenv("OPTSCHED_DEBUG"); + const bool doDebug = dbgEnv && std::atoi(dbgEnv) != 0; - for (auto &MI : DAG) - storedSchedules[schedIndex].push_back(&MI); + // 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); + } + } } // 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 +404,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 +583,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"); @@ -843,6 +878,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 +1028,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")); EnumPriorities = parseHeuristic(schedIni.GetString("ENUM_HEURISTIC")); SecondPassEnumPriorities = parseHeuristic(schedIni.GetString("SECOND_PASS_ENUM_HEURISTIC")); @@ -1019,7 +1060,8 @@ void ScheduleDAGOptSched::loadOptSchedConfig() { OccupancyLimitSource = OCC_LIMIT_TYPE::OLT_NONE; if (ShouldLimitOccupancy) - OccupancyLimitSource = parseOccLimit(schedIni.GetString("OCCUPANCY_LIMIT_SOURCE")); + OccupancyLimitSource = + parseOccLimit(schedIni.GetString("OCCUPANCY_LIMIT_SOURCE")); DeviceACOEnabled = schedIni.GetBool("DEV_ACO"); }