Title: [TACO-26] CPU-Parallel ACO: Persistent Thread Pool and Concurrency Fixes

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

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 <atomic>
+#include <condition_variable>
+#include <functional>
 #include <memory>
+#include <mutex>
+#include <random>
+#include <thread>
+#include <utility>
+#include <vector>
 #include <hip/hip_runtime.h>
 
+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<unsigned, 8> 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<int> &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<std::vector<std::pair<SchedInstruction *, InstCount>>> instScsrs_;
+
+  std::unique_ptr<CPUThreadPool> cpuPool_;
 };
 
 } // namespace opt_sched

include/opt-sched/Scheduler/bb_spill.h

@@ -18,6 +18,8 @@ Last Update:  Apr. 2011
 #include <vector>
 #include <hip/hip_runtime.h>
 
+#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 {

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);
 }
 

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_;