diff --git a/data/rtl-config-vhdl.json b/data/rtl-config-vhdl.json index 3d278943c8..69d5fca62a 100644 --- a/data/rtl-config-vhdl.json +++ b/data/rtl-config-vhdl.json @@ -458,6 +458,16 @@ "name": "handshake.shrsi", "generator": "python $DYNAMATIC/tools/unit-generators/vhdl/vhdl-unit-generator.py -n $MODULE_NAME -o $OUTPUT_DIR/$MODULE_NAME.vhd -t shrsi -p bitwidth=$BITWIDTH extra_signals=$EXTRA_SIGNALS" }, + { + "name": "handshake.init", + "parameters": [ + { + "name": "INITIAL_VALUE", + "type": "unsigned" + } + ], + "generator": "python $DYNAMATIC/tools/unit-generators/vhdl/vhdl-unit-generator.py -n $MODULE_NAME -o $OUTPUT_DIR/$MODULE_NAME.vhd -t init -p bitwidth=$BITWIDTH extra_signals=$EXTRA_SIGNALS initial_value=$INITIAL_VALUE" + }, { "name": "handshake.shrui", "generator": "python $DYNAMATIC/tools/unit-generators/vhdl/vhdl-unit-generator.py -n $MODULE_NAME -o $OUTPUT_DIR/$MODULE_NAME.vhd -t shrui -p bitwidth=$BITWIDTH extra_signals=$EXTRA_SIGNALS" diff --git a/experimental/include/experimental/Support/FtdImplementation.h b/experimental/include/experimental/Support/FtdImplementation.h index 34b68e3abf..b7a2549664 100644 --- a/experimental/include/experimental/Support/FtdImplementation.h +++ b/experimental/include/experimental/Support/FtdImplementation.h @@ -6,9 +6,10 @@ // //===----------------------------------------------------------------------===// // -// Declares the core functions to run the Fast Token Delivery algorithm, -// according to the original FPGA'22 paper by Elakhras et al. -// (https://ieeexplore.ieee.org/document/10035134). +// Declares the top-level steps of the Fast Token Delivery (FTD) algorithm: +// converting phi functions to GSA gates, regeneration, suppression dispatch, +// and condition placeholders. The suppression circuit construction itself +// lives in FtdSuppression.h. // //===----------------------------------------------------------------------===// @@ -18,45 +19,67 @@ #include "dynamatic/Dialect/Handshake/HandshakeOps.h" #include "dynamatic/Support/Backedge.h" #include "experimental/Analysis/GSAAnalysis.h" +#include "experimental/Support/FtdSupport.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" namespace dynamatic { namespace experimental { namespace ftd { +// ShadowCFG is declared in FtdSupport.h + +/// Create SourceOp condition placeholders for every conditional block in the +/// region. Must be called before addGsaGates. +void createAllCondPlaceholders(Region ®ion, OpBuilder &builder); + +/// After addBranchOps (multi-block, handshake ConditionalBranchOps exist), +/// replace each condition placeholder with NotIOp(realCond) that captures +/// the actual handshake condition value. +void resolveCondPlaceholders(handshake::FuncOp funcOp, OpBuilder &builder, + ShadowCFG &shadow); + +/// After addRegen/addSupp, short-circuit all NotIOp condition placeholders +/// and erase them along with their source+constant operands. +void finalizeCondPlaceholders(handshake::FuncOp funcOp); + /// This function implements the regeneration mechanism over a pair made of a /// producer and a consumer (see `addRegen` description). -void addRegenOperandConsumer(PatternRewriter &rewriter, - dynamatic::handshake::FuncOp &funcOp, - Operation *consumerOp, Value operand); +void addRegenOperandConsumer(mlir::OpBuilder &builder, + handshake::FuncOp &funcOp, + mlir::Operation *consumerOp, mlir::Value operand, + ShadowCFG &shadow); /// This function implements the suppression mechanism over a pair made of a /// producer and a consumer (see `addSupp` description). -void addSuppOperandConsumer(PatternRewriter &rewriter, - handshake::FuncOp &funcOp, Operation *consumerOp, - Value operand); +void addSuppOperandConsumer(mlir::OpBuilder &builder, handshake::FuncOp &funcOp, + Operation *consumerOp, Value operand, + ShadowCFG &shadow); /// When the consumer is in a loop while the producer is not, the value must /// be regenerated as many times as needed. This function is in charge of -/// adding some merges to the network, to that this can be done. The new -/// merge is moved inside of the loop, and it works like a reassignment -/// (cfr. FPGA'22, Section V.C). -void addRegen(handshake::FuncOp &funcOp, PatternRewriter &rewriter); +/// adding some merges to the network, so that this can be done. The new +/// merge is moved inside of the loop, and it works like a reassignment. +void addRegen(handshake::FuncOp &funcOp, mlir::OpBuilder &builder, + ShadowCFG &shadow); /// Given each pairs of producers and consumers within the circuit, the /// producer might create a token which is never used by the corresponding /// consumer, because of the control decisions. In this scenario, the token /// must be suppressed. This function inserts a `SUPPRESS` block whenever it -/// is necessary, according to FPGA'22 (IV.C and V) -void addSupp(handshake::FuncOp &funcOp, PatternRewriter &rewriter); +/// is necessary. +void addSupp(handshake::FuncOp &funcOp, mlir::OpBuilder &builder, + ShadowCFG &shadow); /// Starting from the information collected by the gsa analysis pass, /// instantiate some mux operations at the beginning of each block which /// work as explicit phi functions. If `removeTerminators` is true, the `cf` /// terminators in the function are modified to stop feeding the successive /// blocks. -LogicalResult addGsaGates(Region ®ion, PatternRewriter &rewriter, - const gsa::GSAAnalysis &gsa, Backedge startValue, - bool removeTerminators = true); +LogicalResult addGsaGates( + Region ®ion, PatternRewriter &rewriter, const gsa::GSAAnalysis &gsa, + Backedge startValue, + DenseMap> *pendingMuxOperands = nullptr, + bool removeTerminators = true); /// For each non-init merge in the IR, run the GSA analysis to obtain its GSA /// equivalent, then use `addGsaGates` to instantiate such operations in the IR. diff --git a/experimental/include/experimental/Support/FtdSupport.h b/experimental/include/experimental/Support/FtdSupport.h index 0413ac3dad..1ab584fa38 100644 --- a/experimental/include/experimental/Support/FtdSupport.h +++ b/experimental/include/experimental/Support/FtdSupport.h @@ -6,9 +6,10 @@ // //===----------------------------------------------------------------------===// // -// Declares some utility functions which are useful for both the fast token -// delivery algorithm and for the GSA analysis pass. All the functions are about -// analyzing relationships between blocks and handshake operations. +// Declares utility functions and data structures shared across the FTD +// algorithm. Includes CFG analysis helpers, type utilities, annotation +// constants, and the ShadowCFG bridge between the original CFG and the +// flattened handshake IR. // //===----------------------------------------------------------------------===// @@ -18,11 +19,88 @@ #include "dynamatic/Dialect/Handshake/HandshakeOps.h" #include "experimental/Support/BooleanLogic/BoolExpression.h" #include "mlir/Analysis/CFGLoopInfo.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" namespace dynamatic { namespace experimental { namespace ftd { +// ===--------------------------------------------------------------------=== // +// Annotation constants used throughout the FTD algorithm. +// ===--------------------------------------------------------------------=== // + +/// Annotation to use in the IR when an operation needs to be skipped by the FTD +/// algorithm. +constexpr llvm::StringLiteral FTD_OP_TO_SKIP("ftd.skip"); +/// Annotation to identify muxes inserted with the `addGsaGates` +/// functionalities. +constexpr llvm::StringLiteral FTD_EXPLICIT_MU("ftd.MU"); +constexpr llvm::StringLiteral FTD_EXPLICIT_GAMMA("ftd.GAMMA"); +/// Temporary annotation to be used with merges created with the +/// `createPhiNetwork` functionality, which will then be converted into muxes. +constexpr llvm::StringLiteral NEW_PHI("nphi"); +/// Annotation to use for initial merges and initial false constants. +constexpr llvm::StringLiteral FTD_INIT_MERGE("ftd.imerge"); +/// Annotation to use for regeneration multiplexers. +constexpr llvm::StringLiteral FTD_REGEN("ftd.regen"); +/// Annotation for condition variable placeholders. +constexpr llvm::StringLiteral FTD_COND_VAR("ftd.cvar"); + +// ===--------------------------------------------------------------------=== // +// ShadowCFG +// ===--------------------------------------------------------------------=== // + +/// A temporary shadow of the original CFG, built after CfToHandshake +/// conversion flattens everything. Encapsulates the shadow Region +/// (with real CF terminators) plus a condition-value map that bridges +/// shadow analysis to real handshake Values. +/// +/// All analysis infrastructure (BlockIndexing, CFGLoopInfo, path +/// enumeration, dominance) operates on getRegion() as if the original +/// CFG were still alive. The only thing the shadow cannot provide +/// natively is the real handshake condition Value for each cond_br +/// block — getCondition() provides that. +struct ShadowCFG { + mlir::func::FuncOp shadowFunc; + llvm::DenseMap conditionMap; + + mlir::Region &getRegion() { return shadowFunc.getBody(); } + + mlir::Block *getBlock(unsigned bbIdx) { + for (auto [i, blk] : llvm::enumerate(getRegion())) + if (i == bbIdx) + return &blk; + llvm_unreachable("BB index out of range in shadow CFG"); + } + + unsigned getBlockIndex(mlir::Block *block) { + for (auto [i, blk] : llvm::enumerate(getRegion())) + if (&blk == block) + return i; + llvm_unreachable("Block not found in shadow CFG"); + } + + /// Get the real handshake condition Value for the cond_br in block bbIdx. + /// Returns nullptr if the block had an unconditional branch. + mlir::Value getCondition(unsigned bbIdx) { + auto it = conditionMap.find(bbIdx); + return (it != conditionMap.end()) ? it->second : nullptr; + } + + mlir::Value getCondition(mlir::Block *block) { + return getCondition(getBlockIndex(block)); + } + + void destroy() { + if (shadowFunc) + shadowFunc.erase(); + } +}; + +// ===--------------------------------------------------------------------=== // +// BlockIndexing +// ===--------------------------------------------------------------------=== // + /// Class to associate an index to each block, so that if block Bi dominates /// block Bj then i < j. While this is guaranteed by the MLIR CFG construction, /// it cannot really be given for granted, thus it is more convenient to have a @@ -49,10 +127,10 @@ class BlockIndexing { /// Get the index of a block. std::optional getIndexFromBlock(Block *bb) const; - /// Return true if the index of bb1 is greater than then index of bb2. + /// Return true if the index of bb1 is greater than the index of bb2. bool isGreater(Block *bb1, Block *bb2) const; - /// Return true if the index of bb1 is smaller than then index of bb2. + /// Return true if the index of bb1 is smaller than the index of bb2. bool isLess(Block *bb1, Block *bb2) const; /// Given a block whose name is `^BBN` (where N is an integer) return a string @@ -61,6 +139,10 @@ class BlockIndexing { std::string getBlockCondition(Block *block) const; }; +// ===--------------------------------------------------------------------=== // +// CFG Analysis Utilities +// ===--------------------------------------------------------------------=== // + /// Checks if the source and destination are in a loop /// (including any of their ancestor loops). bool isSameLoopBlocks(Block *source, Block *dest, const mlir::CFGLoopInfo &li); @@ -92,12 +174,43 @@ getPathExpression(ArrayRef path, DenseSet &blockIndexSet, const DenseSet &deps = DenseSet(), bool ignoreDeps = true); +/// A lightweight DFS to check if 'end' is reachable from 'start'. +bool isReachable(Block *start, Block *end); + +// ===--------------------------------------------------------------------=== // +// Type Utilities +// ===--------------------------------------------------------------------=== // + /// Return the channelified version of the input type. Type channelifyType(Type type); -/// Get an array of `size` elements all identical to the +/// Get an array of `size` elements all identical to the channelified type. SmallVector getListTypes(Type inputType, unsigned size = 2); +// ===--------------------------------------------------------------------=== // +// IR Attribute Utilities +// ===--------------------------------------------------------------------=== // + +/// Compute the positional index of `block` in its parent region and set +/// the handshake.bb attribute on `op`. +void setBBAttr(Operation *op, Block *block, OpBuilder &builder); + +/// Set the handshake.bb attribute on `op` from an existing attribute. +void setBBAttr(Operation *op, IntegerAttr bbAttr); + +/// Set the handshake.bb attribute on `op`, preferring `bbAttr` if available, +/// otherwise computing from `block`. +void setBBAttrWithFallback(Operation *op, IntegerAttr bbAttr, Block *block, + OpBuilder &builder); + +/// Build a `handshake.bb` IntegerAttr (32-bit unsigned) for `bbIdx`. +IntegerAttr getBBIndexAttr(MLIRContext *ctx, unsigned bbIdx); + +/// Get or create a SourceOp placeholder in `condBlock` representing the +/// condition of that block's terminator. Reuses existing placeholder if one +/// exists. Tagged with FTD_COND_VAR and FTD_OP_TO_SKIP so that FTD skips it. +Value getOrCreateCondPlaceholder(Block *condBlock, OpBuilder &builder); + } // namespace ftd } // namespace experimental } // namespace dynamatic diff --git a/experimental/include/experimental/Support/FtdSuppression.h b/experimental/include/experimental/Support/FtdSuppression.h new file mode 100644 index 0000000000..cdbcf57dac --- /dev/null +++ b/experimental/include/experimental/Support/FtdSuppression.h @@ -0,0 +1,367 @@ +//===- FtdSuppression.h - Suppression Infrastructure ------------*- C++ -*-===// +// +// Dynamatic is under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// Public interface for suppression: building the circuitry that discards a +// producer's data token on the executions where its consumer will not use it. +// +// The entry point is insertDirectSuppression; the remaining declarations are +// the data structures and building blocks it relies on (local CFG and decision +// graph types, loop analysis, Boolean-condition-to-circuit conversion, token +// distribution, and cyclic demotion). See FtdSuppression.cpp for how these fit +// together. +// +//===----------------------------------------------------------------------===// + +#ifndef EXPERIMENTAL_SUPPORT_FTDSUPPRESSION_H +#define EXPERIMENTAL_SUPPORT_FTDSUPPRESSION_H + +#include "dynamatic/Dialect/Handshake/HandshakeOps.h" +#include "dynamatic/Support/Backedge.h" +#include "dynamatic/Support/LLVM.h" +#include "experimental/Support/BooleanLogic/BDD.h" +#include "experimental/Support/BooleanLogic/BoolExpression.h" +#include "experimental/Support/FtdSupport.h" +#include "mlir/Analysis/CFGLoopInfo.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" +#include "mlir/IR/Block.h" +#include "mlir/IR/Dominance.h" +#include "mlir/IR/Region.h" +#include "llvm/ADT/DenseMap.h" +#include "llvm/ADT/DenseSet.h" +#include "llvm/ADT/SmallVector.h" +#include +#include +#include + +namespace dynamatic { +namespace experimental { +namespace ftd { + +// ===--------------------------------------------------------------------=== // +// LocalCFG — Reconstructed subgraph for producer-consumer analysis +// ===--------------------------------------------------------------------=== // + +/// Represents a reconstructed local CFG extracted from the original CFG to +/// represent the producer-consumer relationship. +struct LocalCFG { + // The MLIR region representing the local subgraph. + Region *region = nullptr; + // Mapping: block in local graph -> original block. + DenseMap origMap; + // The producer block in the local CFG. + Block *newProd = nullptr; + // The consumer block in the local CFG. + Block *newCons = nullptr; + // A replicated block used for self-loop delivery (Producer == Consumer). + Block *secondVisitBB = nullptr; + // A unique sink (exit) block to which all terminal paths lead. + Block *sinkBB = nullptr; + // Topological order of the reconstructed region. + SmallVector topoOrder; + // Temporary parent operation that owns the region. + Operation *containerOp = nullptr; + + ~LocalCFG() = default; +}; + +// ===--------------------------------------------------------------------=== // +// LoopScope — Hierarchical loop nesting +// ===--------------------------------------------------------------------=== // + +/// Represents a single level of loop nesting within the local CFG. +/// It forms a tree structure where Level 0 is the top-level acyclic function +/// scope. +struct LoopScope { + /// The nesting level. 0 for the top-level scope, 1 for top-level loops, etc. + unsigned level = 0; + + /// The header block of this loop scope. For Level 0, this is the graph entry. + Block *header = nullptr; + + /// List of latch blocks that jump back to the header in this specific loop. + SmallVector latches; + + /// A set of ALL back-edges contained within this scope, including those + /// belonging to nested sub-loops. Used to prune deep back-edges during DAG + /// extraction. Format: pair. + DenseSet> allBackEdges; + + /// A set of all blocks contained within this scope (including sub-loops). + DenseSet allBlocksInclusive; + + /// List of immediate sub-loops nested within this scope. + SmallVector> subLoops; + + /// Pointer to the parent scope. Null for the top-level scope. + LoopScope *parent = nullptr; + + /// Pointer to the MLIR LoopInfo object. Null for the top-level scope. + mlir::CFGLoop *loopInfo = nullptr; +}; + +// ===--------------------------------------------------------------------=== // +// CyclicGraphManager — Loop analysis and layered CFG extraction +// ===--------------------------------------------------------------------=== // + +/// Manages the cyclic analysis of a LocalCFG. It builds the LoopScope hierarchy +/// and provides utilities to extract acyclic, layered subgraphs for FTD +/// analysis. +class CyclicGraphManager { +public: + /// Constructs the manager and immediately performs topological analysis to + /// build the scope tree. + /// \param lcfg The local control flow graph to analyze. + explicit CyclicGraphManager(LocalCFG &lcfg); + + /// Analyzes the topology of the LocalCFG, identifying loops and building the + /// LoopScope hierarchy (TopLevel -> SubLoops). + void analyzeTopology(); + + /// Returns the nesting level of a given block. + /// \param bb The block to query. + /// \return The nesting level (0 for top-level). + unsigned getNestingLevel(Block *bb) const; + + /// Returns the root of the LoopScope tree (Level 0). + LoopScope *getTopLevelScope() const { return topLevelScope.get(); } + + /// Extracts a normalized, acyclic LocalCFG for a specific LoopScope. + /// This process involves: + /// 1. cloning blocks within the scope. + /// 2. redirecting current-level back-edges to Sink (False). + /// 3. pruning deep-level back-edges (invalid paths). + /// 4. redirecting loop exits to a True terminal. + /// \param scope The loop scope to extract. + /// \param builder The OpBuilder used to create the new graph operations. + /// \return A unique_ptr to the newly created acyclic LocalCFG. + std::unique_ptr extractLayeredCFG(const LoopScope *scope, + OpBuilder &builder); + +private: + /// Reference to the underlying LocalCFG being analyzed. + LocalCFG &lcfg; + + /// Dominator tree analysis utility. + mlir::DominanceInfo domInfo; + + /// Loop analysis utility based on the dominator tree. + mlir::CFGLoopInfo loopInfo; + + /// Mapping from blocks to their nesting level. + DenseMap blockLevelMap; + + /// The root of the scope hierarchy (Level 0). + std::unique_ptr topLevelScope; + + /// Recursively builds the LoopScope tree starting from a given MLIR loop. + /// \param loop The current MLIR loop being analyzed. + /// \param level The current nesting level. + /// \return A unique_ptr to the constructed LoopScope. + std::unique_ptr buildScopeRecursive(mlir::CFGLoop *loop, + unsigned level); +}; + +// ===--------------------------------------------------------------------=== // +// Distribution data structures +// ===--------------------------------------------------------------------=== // + +/// One branch decision along a path: condition variable `var` resolved to +/// `value` (which way the branch went). +struct PathStep { + std::string var; + bool value; + + bool operator==(const PathStep &other) const { + return var == other.var && value == other.value; + } + bool operator!=(const PathStep &other) const { return !(*this == other); } +}; + +/// A path through the control flow, given as the branch decisions taken. +using PathContext = std::vector; + +/// A request for the token of condition variable `varName` on the path `path`. +/// Used while routing condition tokens to the muxes that consume them. +struct VariableRequirement { + std::string varName; + PathContext path; +}; + +/// Maps a variable name to its available versions across different paths. +/// Each entry stores the path context where the value is valid. +struct SignalRegistry { + std::map>> map; + + /// Registers a physical signal available at a specific path context. + void registerSignal(StringRef var, const PathContext &path, Value val); + + /// Finds the best signal source using Longest Prefix Match. + /// Returns the value defined in the deepest matching path context. + Value lookup(StringRef var, const PathContext &queryPath); +}; + +// ===--------------------------------------------------------------------=== // +// CyclicDemotionHelper — Cyclic demotion logic for suppression +// ===--------------------------------------------------------------------=== // + +/// Helper struct encapsulating cyclic demotion logic for reuse across +/// different suppression contexts (direct suppression and MU gate exit). +struct CyclicDemotionHelper { + // Context supplied by the caller (the suppression driver): + mlir::OpBuilder &builder; + MLIRContext *ctx; + const BlockIndexing &bi; + CyclicGraphManager &cyclicMgr; // source of loop scopes and layered CFGs + DenseMap + &origToFullDG; // original block -> decision-graph block + Location loc; + Block *insertBlock; // basic block the emitted ops are tagged to + DenseMap> *pendingMuxOperands; + ShadowCFG *shadow; + // Caches one demoted value per (condition variable, target level) so the + // same demotion circuit is not built twice. + std::map, Value> demotionCache; + + CyclicDemotionHelper( + mlir::OpBuilder &builder, MLIRContext *ctx, const BlockIndexing &bi, + CyclicGraphManager &cyclicMgr, DenseMap &origToFullDG, + Location loc, Block *insertBlock, + DenseMap> *pendingMuxOperands, + ShadowCFG *shadow = nullptr); + + /// Get the nesting level of a condition variable. + unsigned getVarNativeLevel(const std::string &var); + + /// Find the LoopScope at a given level that contains a block. + LoopScope *findScopeForBlock(Block *origIRBlock, unsigned level); + + /// Demotes a value from fromLevel to fromLevel-1 using the layered CFG + /// of the LoopScope at fromLevel that contains origBlock. + Value demoteOneLevel(Value currentValue, Block *origBlock, + unsigned fromLevel); + + /// Returns the value of a condition variable demoted to the target level. + /// Uses caching to avoid duplicate circuits. + Value getValueAtLevel(const std::string &varName, unsigned targetLevel); + + /// Pre-register demoted values for high-level variables in a decision + /// graph into the given registry. + template + void preRegisterDemotedValues(DGType &dg, SignalRegistry ®istry) { + for (Block *b : dg->topoOrder) { + Block *origBlock = dg->origMap.lookup(b); + if (!origBlock) + continue; + std::string var = bi.getBlockCondition(origBlock); + if (var.empty()) + continue; + unsigned native = getVarNativeLevel(var); + if (native > 0) { + Value demoted = getValueAtLevel(var, 0); + if (demoted) + registry.registerSignal(var, {}, demoted); + } + } + } +}; + +// ===--------------------------------------------------------------------=== // +// Graph construction functions +// ===--------------------------------------------------------------------=== // + +/// Build a local control-flow subgraph (LocalCFG) between a producer and +/// consumer. The subgraph is reconstructed as a region with unique entry +/// (producer) and exit (sink). +std::unique_ptr buildLocalCFGRegion(OpBuilder &builder, + Block *origProd, Block *origCons, + const BlockIndexing &bi); + +/// Constructs a NEW LocalCFG that represents the Decision Graph from a +/// raw LocalCFG, retaining only the blocks in `dependencies` plus the +/// consumer and sink. +std::unique_ptr buildDecisionGraph( + const LocalCFG &rawGraph, const DenseSet &dependencies, + const DenseMap &muxConstraints = DenseMap()); + +// ===--------------------------------------------------------------------=== // +// Path enumeration and expression generation +// ===--------------------------------------------------------------------=== // + +/// Enumerates all paths from producer to consumer in a LocalCFG and returns +/// the minimized SOP boolean expression representing the reaching condition. +boolean::BoolExpression *enumeratePaths(const LocalCFG &lcfg, + const BlockIndexing &bi, + const DenseSet &controlDeps); + +// ===--------------------------------------------------------------------=== // +// Circuit construction (BDD → hardware) +// ===--------------------------------------------------------------------=== // + +/// Recursively converts a BDD to a Mux Tree circuit. +Value bddToCircuit( + mlir::OpBuilder &builder, boolean::BDD *bdd, Block *block, + SignalRegistry ®istry, PathContext currentPath, const BlockIndexing &bi, + DenseMap> *pendingMuxOperands, + ShadowCFG *shadow = nullptr, IntegerAttr forcedBBAttr = {}); + +/// Main entry point of the token distribution logic. +void buildDistributionNetwork( + mlir::OpBuilder &builder, const LocalCFG &lcfg, const BlockIndexing &bi, + SignalRegistry ®istry, + DenseMap> *pendingMuxOperands, + ShadowCFG *shadow = nullptr); + +/// Compute the topological rank of original blocks from a LocalCFG. +/// Maps each original block (via origMap) to its index in topoOrder. +DenseMap computeTopoRank(const LocalCFG &lcfg); + +/// Convert a boolean expression to a hardware mux-tree circuit. +/// Handles the full pipeline: minimize → topo-sort cofactors → buildBDD → +/// bddToCircuit. `varRank` maps original CFG blocks to their topological rank +/// (determines BDD variable ordering). +Value expressionToCircuit( + OpBuilder &builder, boolean::BoolExpression *expr, + const DenseMap &varRank, Block *insertBlock, + SignalRegistry ®istry, const BlockIndexing &bi, + DenseMap> *pendingMuxOperands = nullptr, + ShadowCFG *shadow = nullptr, IntegerAttr forcedBBAttr = {}); + +/// Convenience overload: computes varRank from a LocalCFG's topoOrder. +Value expressionToCircuit( + OpBuilder &builder, boolean::BoolExpression *expr, + const LocalCFG &rankSource, Block *insertBlock, SignalRegistry ®istry, + const BlockIndexing &bi, + DenseMap> *pendingMuxOperands = nullptr, + ShadowCFG *shadow = nullptr, IntegerAttr forcedBBAttr = {}); + +/// Compute the loop backedge condition for a self-loop at loopHeader. +/// Runs the full suppression pipeline for a header-to-header path: +/// buildLocalCFGRegion → CDA → CyclicGraphManager → CyclicDemotionHelper → +/// buildDistributionNetwork → enumeratePaths → expressionToCircuit. +/// Returns the condition Value that is true when the loop iterates back. +Value computeLoopBackedgeCondition( + OpBuilder &builder, Block *loopHeader, Block *insertBlock, + const BlockIndexing &bi, + DenseMap> *pendingMuxOperands = nullptr, + ShadowCFG *shadow = nullptr); + +// ===--------------------------------------------------------------------=== // +// Top-level suppression entry point +// ===--------------------------------------------------------------------=== // + +/// Apply the suppression algorithm for a non-loop producer-consumer pair. +/// This is the main entry point for building suppression circuits. +void insertDirectSuppression(mlir::OpBuilder &builder, + handshake::FuncOp &funcOp, Operation *consumer, + Value connection, ShadowCFG &shadow); + +} // namespace ftd +} // namespace experimental +} // namespace dynamatic + +#endif // EXPERIMENTAL_SUPPORT_FTDSUPPRESSION_H diff --git a/experimental/lib/Analysis/GSAAnalysis.cpp b/experimental/lib/Analysis/GSAAnalysis.cpp index b8fa890117..ec5d12f706 100644 --- a/experimental/lib/Analysis/GSAAnalysis.cpp +++ b/experimental/lib/Analysis/GSAAnalysis.cpp @@ -236,16 +236,7 @@ experimental::gsa::Gate *experimental::gsa::GSAAnalysis::expandGammaTree( ++uniqueGateIndex, bi.getBlockFromIndex(indexToUse).value(), BoolExpression::boolVar(conditionToUse), {conditionToUse}); // since condition is one block boolvar is enough - - // If the Gamma is a result of the expansion of a Mu that has more than two - // inputs, force its placement in the block of its condition because placing - // it in the block of the Mu, which is always a loop header, will mess up the - // control dependence analysis betweem the newly inserted Gamma and its - // producers that are in the loop body in this case - if (originalPhi->muGenerated) - newGate->gateBlock = newGate->conditionBlock; - - gatesPerBlock[newGate->getBlock()].push_back(newGate); + gatesPerBlock[originalPhi->getBlock()].push_back(newGate); return newGate; } diff --git a/experimental/lib/Conversion/FtdCfToHandshake.cpp b/experimental/lib/Conversion/FtdCfToHandshake.cpp index 5b0864b7e6..cb347fa1a8 100644 --- a/experimental/lib/Conversion/FtdCfToHandshake.cpp +++ b/experimental/lib/Conversion/FtdCfToHandshake.cpp @@ -13,21 +13,30 @@ #include "dynamatic/Dialect/Handshake/HandshakeDialect.h" #include "dynamatic/Dialect/Handshake/HandshakeInterfaces.h" #include "dynamatic/Dialect/Handshake/HandshakeOps.h" +#include "dynamatic/Dialect/Handshake/HandshakeTypes.h" +#include "dynamatic/Support/Backedge.h" #include "dynamatic/Support/CFG.h" #include "experimental/Support/CFGAnnotation.h" #include "experimental/Support/FtdImplementation.h" +#include "mlir/Analysis/CFGLoopInfo.h" #include "mlir/Dialect/Affine/Utils.h" #include "mlir/Dialect/Arith/IR/Arith.h" #include "mlir/Dialect/ControlFlow/IR/ControlFlow.h" +#include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h" #include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/Math/IR/Math.h" +#include "mlir/Dialect/MemRef/IR/MemRef.h" #include "mlir/IR/BuiltinAttributes.h" #include "mlir/IR/BuiltinDialect.h" #include "mlir/IR/BuiltinOps.h" +#include "mlir/IR/Dominance.h" #include "mlir/IR/MLIRContext.h" #include "mlir/IR/PatternMatch.h" #include "mlir/Support/LogicalResult.h" #include "mlir/Transforms/DialectConversion.h" +#include "llvm/ADT/SmallVector.h" + +#include #include // [START Boilerplate code for the MLIR pass] @@ -46,6 +55,216 @@ using namespace dynamatic::experimental; using namespace dynamatic::experimental::boolean; using namespace dynamatic::experimental::ftd; +struct AllocaOpConversion : public DynOpConversionPattern { + using DynOpConversionPattern::DynOpConversionPattern; + + // Construct a dense element attribute with everything zeroes. + DenseElementsAttr getZeroAttr(ShapedType type) const { + auto elemType = type.getElementType(); + if (auto intTy = dyn_cast(elemType)) { + return DenseElementsAttr::get(type, APInt(intTy.getWidth(), 0)); + } + if (auto floatTy = dyn_cast(type)) { + if (floatTy.isF16()) + return DenseElementsAttr::get( + type, APFloat::getZero(APFloat::IEEEhalf(), /*negative=*/false)); + if (floatTy.isBF16()) + return DenseElementsAttr::get( + type, APFloat::getZero(APFloat::BFloat(), /*negative=*/false)); + if (floatTy.isF32()) + return DenseElementsAttr::get( + type, APFloat::getZero(APFloat::IEEEsingle(), /*negative=*/false)); + if (floatTy.isF64()) + return DenseElementsAttr::get( + type, APFloat::getZero(APFloat::IEEEdouble(), /*negative=*/false)); + llvm::report_fatal_error("Unhandled float element type!"); + } + llvm::report_fatal_error("Unknown base element type!"); + } + + LogicalResult + matchAndRewrite(memref::AllocaOp op, OpAdaptor adapter, + ConversionPatternRewriter &rewriter) const override { + // HACK: By default, we initialize the memory with all zeros. According to + // the C standard, this only happens for arrays. + rewriter.replaceOpWithNewOp(op, op.getType(), + getZeroAttr(op.getType())); + return success(); + } +}; + +struct GetGlobalOpConversion + : public DynOpConversionPattern { + using DynOpConversionPattern::DynOpConversionPattern; + LogicalResult + matchAndRewrite(memref::GetGlobalOp op, OpAdaptor adapter, + ConversionPatternRewriter &rewriter) const override { + // clang-format off + // Example: + // memref.global "external" constant @internal_array : memref<...> = dense<...> + // .... + // %4 = memref.get_global @internal_array : memref<...> + // + // In this case, we remove the global constant and rewrite the addressof + // node into a RAMOp (and we put an attribute to describe its constant + // value). + // clang-format on + SymbolTableCollection symbolTableCollection; + + auto symNameOfGetGlobal = op.getNameAttr(); + + memref::GlobalOp global; + auto moduleOp = op->getParentOfType(); + moduleOp.walk([&global, symNameOfGetGlobal](memref::GlobalOp gbl) { + if (gbl.getSymName() == symNameOfGetGlobal.getValue()) { + global = gbl; + } + }); + + if (!global) { + // No corresponding Global (maybe emit pass failure is better) + return failure(); + } + + /// The initial value doesn't have any type constraints. Therefore we need + /// to check if it is stored as dense elements. + mlir::Attribute initValueAttr = global.getInitialValueAttr(); + if (auto denseAttr = initValueAttr.dyn_cast()) { + rewriter.replaceOpWithNewOp(op, op.getType(), + denseAttr); + } else { + llvm::report_fatal_error( + "The initial value must be denoted in DenseElementsAttr."); + } + return success(); + } +}; + +// TODO: Here we simply erase all the global variables and attach the initial +// values to the RAMOps inside the handshake function. +struct GlobalOpConversion : public DynOpConversionPattern { + using DynOpConversionPattern::DynOpConversionPattern; + LogicalResult + matchAndRewrite(memref::GlobalOp op, OpAdaptor adapter, + ConversionPatternRewriter &rewriter) const override { + rewriter.eraseOp(op); + return success(); + } +}; + +/// Per-block edge information captured from CF-level IR before conversion. +struct BlockEdgeInfo { + bool isConditional = false; + bool hasSuccessors = false; + unsigned trueSuccIdx = 0; + unsigned falseSuccIdx = 0; + unsigned uncondSuccIdx = 0; +}; + +/// Complete CFG topology of one function, captured before conversion. +struct OriginalCFGInfo { + unsigned numBlocks = 0; + SmallVector blockEdges; +}; + +/// Walk every func::FuncOp in the module and capture its CFG topology. +/// Must be called BEFORE applyFullConversion. +static DenseMap +captureAllCFGTopologies(ModuleOp moduleOp) { + DenseMap result; + + for (auto funcOp : moduleOp.getOps()) { + if (funcOp.isExternal() || funcOp.getSymName().startswith("__init")) + continue; + + Region ®ion = funcOp.getBody(); + if (region.empty()) + continue; + + OriginalCFGInfo info; + + DenseMap blockIdx; + for (auto [idx, block] : llvm::enumerate(region)) + blockIdx[&block] = idx; + + info.numBlocks = blockIdx.size(); + info.blockEdges.resize(info.numBlocks); + + for (auto &[block, idx] : blockIdx) { + BlockEdgeInfo &edge = info.blockEdges[idx]; + Operation *term = block->getTerminator(); + + if (auto condBr = dyn_cast(term)) { + edge.isConditional = true; + edge.hasSuccessors = true; + edge.trueSuccIdx = blockIdx.lookup(condBr.getTrueDest()); + edge.falseSuccIdx = blockIdx.lookup(condBr.getFalseDest()); + } else if (auto br = dyn_cast(term)) { + edge.hasSuccessors = true; + edge.uncondSuccIdx = blockIdx.lookup(br.getDest()); + } + } + + result[funcOp.getSymName()] = std::move(info); + } + + return result; +} + +/// Create a ShadowCFG for a given handshake::FuncOp using the topology +/// captured before conversion. +static ftd::ShadowCFG buildShadowCFG(OpBuilder &builder, + handshake::FuncOp realFuncOp, + const OriginalCFGInfo &info) { + ftd::ShadowCFG shadow; + Location loc = realFuncOp.getLoc(); + + // 1. Create a temporary func::FuncOp with blocks + CF terminators + { + OpBuilder::InsertionGuard guard(builder); + builder.setInsertionPointAfter(realFuncOp); + auto funcType = builder.getFunctionType({}, {}); + shadow.shadowFunc = + builder.create(loc, "__ftd_shadow_cfg__", funcType); + + Region &R = shadow.shadowFunc.getBody(); + SmallVector blocks; + for (unsigned i = 0; i < info.numBlocks; ++i) + blocks.push_back(builder.createBlock(&R, R.end())); + + for (unsigned i = 0; i < info.numBlocks; ++i) { + const BlockEdgeInfo &edge = info.blockEdges[i]; + builder.setInsertionPointToEnd(blocks[i]); + + if (edge.isConditional) { + auto dummyCond = + builder.create(loc, builder.getBoolAttr(true)); + builder.create( + loc, dummyCond, blocks[edge.trueSuccIdx], ValueRange{}, + blocks[edge.falseSuccIdx], ValueRange{}); + } else if (edge.hasSuccessors) { + builder.create(loc, blocks[edge.uncondSuccIdx]); + } else { + builder.create(loc); + } + } + } + + // 2. Scan the real funcOp to map BB index -> real condition Value + realFuncOp.walk([&](handshake::ConditionalBranchOp brOp) { + if (brOp->hasAttr("ftd.skip")) + return; + auto bbAttr = brOp->getAttrOfType("handshake.bb"); + if (!bbAttr) + return; + unsigned bbIdx = bbAttr.getUInt(); + if (!shadow.conditionMap.contains(bbIdx)) + shadow.conditionMap[bbIdx] = brOp.getConditionOperand(); + }); + + return shadow; +} + namespace { struct FtdCfToHandshakePass @@ -59,45 +278,53 @@ struct FtdCfToHandshakePass CfToHandshakeTypeConverter converter; RewritePatternSet patterns(ctx); + // Capture CFG topology before conversion flattens everything. + auto cfgTopologies = captureAllCFGTopologies(modOp); + patterns.add( getAnalysis(), getAnalysis(), getAnalysis(), converter, ctx); - patterns - .add, - FtdConvertIndexCast, - FtdOneToOneConversion, - FtdOneToOneConversion, - FtdOneToOneConversion, - FtdOneToOneConversion, - FtdOneToOneConversion, - FtdOneToOneConversion, - FtdOneToOneConversion, - FtdOneToOneConversion, - FtdOneToOneConversion, - FtdOneToOneConversion, - FtdOneToOneConversion, - FtdOneToOneConversion, - FtdOneToOneConversion, - FtdOneToOneConversion, - FtdOneToOneConversion, - FtdOneToOneConversion, - FtdOneToOneConversion, - FtdOneToOneConversion, - FtdOneToOneConversion, - FtdOneToOneConversion, - FtdOneToOneConversion, - FtdOneToOneConversion, - FtdOneToOneConversion, - FtdOneToOneConversion, - FtdOneToOneConversion, - FtdOneToOneConversion, - FtdOneToOneConversion, - FtdOneToOneConversion, - FtdOneToOneConversion>( - getAnalysis(), converter, ctx); + patterns.add< + // LowerFuncToHandshake, + /*ConvertConstants,*/ AllocaOpConversion, ConvertCalls, + /*ConvertUndefinedValues,*/ GetGlobalOpConversion, GlobalOpConversion, + ConvertIndexCast, + ConvertIndexCast, + OneToOneConversion, + OneToOneConversion, + OneToOneConversion, + OneToOneConversion, + OneToOneConversion, + OneToOneConversion, + OneToOneConversion, + OneToOneConversion, + OneToOneConversion, + OneToOneConversion, + OneToOneConversion, + OneToOneConversion, + OneToOneConversion, + OneToOneConversion, + OneToOneConversion, + OneToOneConversion, + OneToOneConversion, + OneToOneConversion, + OneToOneConversion, + OneToOneConversion, + OneToOneConversion, + OneToOneConversion, + OneToOneConversion, + OneToOneConversion, + OneToOneConversion, + OneToOneConversion, + OneToOneConversion, + OneToOneConversion, + OneToOneConversion, + OneToOneConversion, + OneToOneConversion, + OneToOneConversion>( + getAnalysis(), converter, ctx); // All func-level functions must become handshake-level functions ConversionTarget target(*ctx); @@ -107,28 +334,94 @@ struct FtdCfToHandshakePass arith::ArithDialect, math::MathDialect, BuiltinDialect>(); + target.addDynamicallyLegalOp([](func::CallOp op) { + // If the call is to __init, consider it legal for now. + // This allows the pass to continue so that __placeholder conversion can + // later erase these __init calls. + // All other func.CallOp (not calling __init) remain illegal due to the + // addIllegalDialect rule above and must be converted by a pattern. + if (auto calledFn = dyn_cast_or_null( + SymbolTable::lookupNearestSymbolFrom(op, op.getCalleeAttr()))) { + return calledFn.getSymName().startswith("__init"); + } + // If symbol lookup fails or it's not a func::FuncOp, treat as default + // (illegal) + return false; + }); + target.addDynamicallyLegalOp( + [](func::FuncOp op) { return op.getSymName().startswith("__init"); }); + if (failed(applyFullConversion(modOp, target, std::move(patterns)))) return signalPassFailure(); - } -}; -} // namespace -using ArgReplacements = DenseMap; + // Clean up: Remove the definition of each __init* function, but only if it + // has no remaining uses. This is safe because all valid calls to __init* + // were tracked and deleted earlier. + for (auto func : llvm::make_early_inc_range(modOp.getOps())) { + if (func.getSymName().startswith("__init")) { + assert(func.use_empty() && + "__init function should not have users after transformation"); + func.erase(); + } + } -static void channelifyMuxes(handshake::FuncOp &funcOp) { - // Considering each mux that was added, the inputs and output values must be - // channellified - for (handshake::MuxOp muxOp : funcOp.getOps()) { - assert(muxOp.getDataOperands().size() == 2 && - "Multiplexers should have two data inputs"); - muxOp.getDataOperands()[0].setType( - channelifyType(muxOp.getDataOperands()[0].getType())); - muxOp.getDataOperands()[1].setType( - channelifyType(muxOp.getDataOperands()[1].getType())); - muxOp.getDataResult().setType( - channelifyType(muxOp.getDataResult().getType())); + for (auto funcOp : modOp.getOps()) { + mlir::OpBuilder builder(funcOp.getContext()); + + auto topoIt = cfgTopologies.find(funcOp.getName()); + if (topoIt == cfgTopologies.end()) + continue; + const OriginalCFGInfo &info = topoIt->second; + + if (info.numBlocks <= 1) + continue; + + // Build the shadow CFG — one struct, everything inside. + ftd::ShadowCFG shadow = buildShadowCFG(builder, funcOp, info); + + // Route the select of every conditional block's terminator branch + // through that block's condition placeholder. + DenseMap condPlaceholderByBB; + for (auto cstOp : funcOp.getOps()) { + if (!cstOp->hasAttr("ftd.cvar")) + continue; + if (auto bbAttr = cstOp->getAttrOfType("handshake.bb")) + condPlaceholderByBB[bbAttr.getUInt()] = cstOp.getResult(); + } + for (auto brOp : funcOp.getOps()) { + if (brOp->hasAttr("ftd.skip")) + continue; + auto bbAttr = brOp->getAttrOfType("handshake.bb"); + if (!bbAttr) + continue; + auto it = condPlaceholderByBB.find(bbAttr.getUInt()); + if (it == condPlaceholderByBB.end()) + continue; + // operand 0 of a ConditionalBranchOp is the condition (select). + brOp->setOperand(0, it->second); + } + + ftd::resolveCondPlaceholders(funcOp, builder, shadow); + + // Populate conditionMap from NotIOp placeholders + for (auto notOp : funcOp.getOps()) { + if (!notOp->hasAttr("ftd.cvar")) + continue; + auto bbAttr = notOp->getAttrOfType("handshake.bb"); + if (!bbAttr) + continue; + shadow.conditionMap[bbAttr.getUInt()] = notOp.getResult(); + } + + ftd::addRegen(funcOp, builder, shadow); + ftd::addSupp(funcOp, builder, shadow); + ftd::finalizeCondPlaceholders(funcOp); + + shadow.destroy(); + } } -} +}; +} // namespace /// Converts undefined operations (LLVM::UndefOp) with a default "0" /// constant triggered by the start signal of the corresponding function. @@ -172,6 +465,27 @@ static LogicalResult convertUndefinedValues(ConversionPatternRewriter &rewriter, return success(); } +/// Determines whether it is possible to transform an arith-level constant into +/// a Handshake-level constant that is triggered by an always-triggering source +/// component without compromising the circuit semantics (e.g., without +/// triggering a memory operation before the circuit "starts"). Returns false if +/// the Handshake-level constant that replaces the input must instead be +/// connected to the control-only network; returns true otherwise. This function +/// assumes that the rest of the std-level operations have already been +/// converted to their Handshake equivalent. +/// NOTE: I doubt this works in half-degenerate cases, but this is the logic +/// that legacy Dynamatic follows. +static bool isCstSourcable(arith::ConstantOp cstOp) { + std::function isValidUser = [&](Operation *user) -> bool { + if (isa(user)) + return llvm::all_of(user->getUsers(), isValidUser); + return !isa(user); + }; + + return llvm::all_of(cstOp->getUsers(), isValidUser); +} + /// Convers arith-level constants to handshake-level constants. Constants are /// triggered by the start value of the corresponding function. The FTD /// algorithm is then in charge of connecting the constants to the rest of the @@ -193,7 +507,14 @@ static LogicalResult convertConstants(ConversionPatternRewriter &rewriter, // This variable will work as activation value for the constant. If the // constant is considered as sourcable, this will be the output of a source // component, otherwise it remains startValue - auto controlValue = startValue; + Value controlValue; + if (isCstSourcable(cstOp)) { + auto sourceOp = rewriter.create(cstOp.getLoc()); + inheritBB(cstOp, sourceOp); + controlValue = sourceOp.getResult(); + } else { + controlValue = startValue; + } // Continue the conversion by obtaining the size of the constnat TypedAttr valueAttr = cstOp.getValue(); @@ -217,34 +538,10 @@ static LogicalResult convertConstants(ConversionPatternRewriter &rewriter, return success(); } -template -LogicalResult FtdOneToOneConversion::matchAndRewrite( - SrcOp srcOp, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const { - rewriter.setInsertionPoint(srcOp); - SmallVector newTypes; - for (Type resType : srcOp->getResultTypes()) - newTypes.push_back(channelifyType(resType)); - auto newOp = - rewriter.create(srcOp->getLoc(), newTypes, adaptor.getOperands(), - srcOp->getAttrDictionary().getValue()); - - // /!\ This is the main difference from the base function. Without such - // replacement, a "null operand found" error is present at the end of the - // transformation pass in almost any test. This is due to the way FTD tweaks - // the coexistence of `cf` and `handshake` dialect to obtain a final circuit: - // without such explicit replacement, deleted operations still provide values - // to new operations. However, this should be fixed by understanding what is - // causing MLIR to complain. - for (auto [from, to] : llvm::zip(srcOp->getResults(), newOp->getResults())) - from.replaceAllUsesWith(to); - - this->namer.replaceOp(srcOp, newOp); - rewriter.replaceOp(srcOp, newOp); - return success(); -} +using ArgReplacements = DenseMap; LogicalResult ftd::FtdLowerFuncToHandshake::matchAndRewrite( - func::FuncOp lowerFuncOp, OpAdaptor adaptor, + func::FuncOp lowerFuncOp, OpAdaptor /*adaptor*/, ConversionPatternRewriter &rewriter) const { // Map all memory accesses in the matched function to the index of their // memref in the function's arguments @@ -254,14 +551,22 @@ LogicalResult ftd::FtdLowerFuncToHandshake::matchAndRewrite( memrefToArgIdx.insert({arg, idx}); } - // Add the muxes as obtained by the GSA analysis pass. This requires the start - // value, as init merges need it as one of their output. However, the start - // value is not available yet here, so a backedge is adopted instead. + ftd::createAllCondPlaceholders(lowerFuncOp.getRegion(), rewriter); + + // Structure used inside addGsaGates to temporarily map a cf value to a + // backedge until the proper handshake values are created; in which case, the + // backedge is replaced with the corresponding hanshake values + static DenseMap> pendingMuxOperands; + + // Add the muxes as obtained by the GSA analysis pass. This requires the + // start value, as init merges need it as one of their output. However, + // the start value is not available yet here, so a backedge is adopted + // instead. BackedgeBuilder edgeBuilderStart(rewriter, lowerFuncOp.getRegion().getLoc()); Backedge startValueBackedge = edgeBuilderStart.get(rewriter.getType()); if (failed(addGsaGates(lowerFuncOp.getRegion(), rewriter, gsaAnalysis, - startValueBackedge))) + startValueBackedge, &pendingMuxOperands))) return failure(); // First lower the parent function itself, without modifying its body @@ -279,6 +584,15 @@ LogicalResult ftd::FtdLowerFuncToHandshake::matchAndRewrite( // it. startValueBackedge.setValue((Value)funcOp.getArguments().back()); + for (auto &[originalValue, backedges] : pendingMuxOperands) { + Value newVal = rewriter.getRemappedValue(originalValue); + assert(newVal && "Failed to remap GSA mux operand!"); + + for (Backedge &be : backedges) + be.setValue(newVal); + } + pendingMuxOperands.clear(); + // Stores mapping from each value that passes through a merge-like // operation to the data result of that merge operation ArgReplacements argReplacements; @@ -292,12 +606,36 @@ LogicalResult ftd::FtdLowerFuncToHandshake::matchAndRewrite( addMergeOps(funcOp, rewriter, argReplacements); addBranchOps(funcOp, rewriter); + // addBranchOps only creates handshake::ConditionalBranchOp for live-out + // values. If a conditional block has no live-outs, no ConditionalBranchOp + // is created and the condition value is lost after flattening. + // Create one using the block's control signal so that buildShadowCFG's + // walk can always find the condition. + for (Block &block : funcOp) { + auto condBr = dyn_cast(block.getTerminator()); + if (!condBr) + continue; + + bool hasHandshakeCondBr = llvm::any_of(block, [](Operation &op) { + return isa(&op); + }); + + if (!hasHandshakeCondBr) { + Value cond = rewriter.getRemappedValue(condBr.getCondition()); + assert(cond && "Failed to remap condition"); + Value ctrl = block.getArguments().back(); + rewriter.setInsertionPoint(condBr); + rewriter.create(condBr.getLoc(), cond, + ctrl); + } + } + // The memory operations are converted to the corresponding handshake // counterparts. No LSQ interface is created yet. BackedgeBuilder edgeBuilder(rewriter, funcOp->getLoc()); LowerFuncToHandshake::MemInterfacesInfo memInfo; if (failed(convertMemoryOps(funcOp, rewriter, memrefToArgIdx, edgeBuilder, - memInfo, true))) + memInfo))) return failure(); // First round of bb-tagging so that newly inserted Dynamatic memory ports @@ -314,20 +652,10 @@ LogicalResult ftd::FtdLowerFuncToHandshake::matchAndRewrite( // Convert the constants and undefined values from the `arith` dialect to // the `handshake` dialect, while also using the start value as their // control value - if (failed(::convertConstants(rewriter, funcOp, namer)) || - failed(::convertUndefinedValues(rewriter, funcOp, namer))) + if (failed(convertConstants(rewriter, funcOp, namer)) || + failed(convertUndefinedValues(rewriter, funcOp, namer))) return failure(); - if (funcOp.getBlocks().size() != 1) { - - // Add muxes for regeneration of values in loop - addRegen(funcOp, rewriter); - channelifyMuxes(funcOp); - - // Add suppression blocks between each pair of producer and consumer - addSupp(funcOp, rewriter); - } - // id basic block idBasicBlocks(funcOp, rewriter); @@ -339,47 +667,3 @@ LogicalResult ftd::FtdLowerFuncToHandshake::matchAndRewrite( return success(); } - -template -LogicalResult FtdConvertIndexCast::matchAndRewrite( - CastOp castOp, OpAdaptor adaptor, - ConversionPatternRewriter &rewriter) const { - - auto getWidth = [](Type type) -> unsigned { - // In Fast Token Delivery the type of the element might be already a - // channel, rather than a simple type. In this case, the type should be - // extracted. We also make sure that no extra bits are present at this - // compilation stage. - if (auto dataType = dyn_cast(type)) { - assert(dataType.getNumExtraSignals() == 0 && - "expected type to have no extra signals"); - type = dataType.getDataType(); - } - if (isa(type)) - return 32; - return type.getIntOrFloatBitWidth(); - }; - - unsigned srcWidth = getWidth(castOp.getOperand().getType()); - unsigned dstWidth = getWidth(castOp.getResult().getType()); - Type dstType = handshake::ChannelType::get(rewriter.getIntegerType(dstWidth)); - Operation *newOp; - if (srcWidth < dstWidth) { - // This is an extension - newOp = - rewriter.create(castOp.getLoc(), dstType, adaptor.getOperands(), - castOp->getAttrDictionary().getValue()); - } else { - // This is a truncation - newOp = rewriter.create( - castOp.getLoc(), dstType, adaptor.getOperands(), - castOp->getAttrDictionary().getValue()); - } - this->namer.replaceOp(castOp, newOp); - rewriter.replaceOp(castOp, newOp); - - // /!\ This is again the main difference from the normal flow. See the comment - // in FtdOneToOneConversion. - castOp.getResult().replaceAllUsesWith(newOp->getResult(0)); - return success(); -} diff --git a/experimental/lib/Support/CFGAnnotation.cpp b/experimental/lib/Support/CFGAnnotation.cpp index 0502445e54..55c2f130d3 100644 --- a/experimental/lib/Support/CFGAnnotation.cpp +++ b/experimental/lib/Support/CFGAnnotation.cpp @@ -270,7 +270,14 @@ dynamatic::experimental::cfg::restoreCfStructure(handshake::FuncOp &funcOp, for (auto &[source, _] : edges) blocksList.push_back(source); std::sort(blocksList.begin(), blocksList.end()); - blocksList.push_back(blocksList.back() + 1); + if (blocksList.empty()) { + // A function with a single block and no outgoing branches serializes to an + // empty edge set. Rebuild it as the trivial block 0 CFG instead of + // dereferencing an empty vector. + blocksList.push_back(0); + } else { + blocksList.push_back(blocksList.back() + 1); + } // Maintains the current block under analysis: all the operations in block 0 // are maintained in the same location, while the others operations are diff --git a/experimental/lib/Support/CMakeLists.txt b/experimental/lib/Support/CMakeLists.txt index 5433a4a9f1..1c86dfe55e 100644 --- a/experimental/lib/Support/CMakeLists.txt +++ b/experimental/lib/Support/CMakeLists.txt @@ -7,6 +7,7 @@ add_dynamatic_library(DynamaticExperimentalSupport HandshakeSimulator.cpp FtdImplementation.cpp FtdSupport.cpp + FtdSuppression.cpp CFGAnnotation.cpp FormalProperty.cpp FlowExpression.cpp diff --git a/experimental/lib/Support/FtdImplementation.cpp b/experimental/lib/Support/FtdImplementation.cpp index 9ab2726c26..f57fe9c6bc 100644 --- a/experimental/lib/Support/FtdImplementation.cpp +++ b/experimental/lib/Support/FtdImplementation.cpp @@ -6,9 +6,9 @@ // //===----------------------------------------------------------------------===// // -// Implements the core functions to run the Fast Token Delivery algorithm, -// according to the original FPGA'22 paper by Elakhras et al. -// (https://ieeexplore.ieee.org/document/10035134). +// Implements the top-level FTD algorithm orchestration: GSA conversion, +// regeneration, suppression dispatch, and phi networks. The suppression +// circuit construction infrastructure lives in FtdSuppression.cpp. // //===----------------------------------------------------------------------===// @@ -18,6 +18,7 @@ #include "experimental/Support/BooleanLogic/BDD.h" #include "experimental/Support/BooleanLogic/BoolExpression.h" #include "experimental/Support/FtdSupport.h" +#include "experimental/Support/FtdSuppression.h" #include "mlir/Analysis/CFGLoopInfo.h" #include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" @@ -26,63 +27,85 @@ using namespace mlir; using namespace dynamatic; using namespace dynamatic::experimental; +using namespace dynamatic::experimental::ftd; using namespace dynamatic::experimental::boolean; -/// Different types of loop suppression. -enum BranchToLoopType { - - // In this case, the producer is inside a loop, while the consumer is outside. - // The token must be suppressed as long as the loop is executed, in order to - // provide only the final token handled. - MoreProducerThanConsumers, - - // In this case, the producer is the consumer itself; this is the case of a - // regeneration multiplexer. The token must be suppressed only if the loop is - // done iterating. - SelfRegeneration, - - // In this case, the token is used back in a loop. The token is to be - // suppressed only if the loop is done iterating. - BackwardRelationship -}; - -/// Annotation to use in the IR when an operation needs to be skipped by the FTD -/// algorithm. -constexpr llvm::StringLiteral FTD_OP_TO_SKIP("ftd.skip"); -/// Annotation to use when a suppression branch is added which needs to go -/// through the suppression mechanism again. -constexpr llvm::StringLiteral FTD_NEW_SUPP("ftd.supp"); -/// Annotation to to identify muxes inserted with the `addGsaGates` -/// functionalities. -constexpr llvm::StringLiteral FTD_EXPLICIT_MU("ftd.MU"); -constexpr llvm::StringLiteral FTD_EXPLICIT_GAMMA("ftd.GAMMA"); -/// Temporary annotation to be used with merges created with the -/// `createPhiNetwork` functionality, which will then be converted into muxes. -constexpr llvm::StringLiteral NEW_PHI("nphi"); -/// Annotation to use for initial merges and initial false constants. -constexpr llvm::StringLiteral FTD_INIT_MERGE("ftd.imerge"); -/// Annotation to use for regeneration multiplexers. -constexpr llvm::StringLiteral FTD_REGEN("ftd.regen"); - -/// Identify the block that has muxCondition as its terminator condition -/// Note that it is not necessarily the same block defining the muxCondition -static Block *returnMuxConditionBlock(Value muxCondition) { - Block *muxConditionBlock = nullptr; - - for (auto &use : muxCondition.getUses()) { - Operation *userOp = use.getOwner(); - Block *userBlock = userOp->getBlock(); - - if (isa_and_nonnull(userOp)) { - muxConditionBlock = userBlock; - break; - } +// ===--------------------------------------------------------------------=== // +// Condition placeholder management +// ===--------------------------------------------------------------------=== // + +void ftd::createAllCondPlaceholders(Region ®ion, OpBuilder &builder) { + for (Block &block : region) { + if (isa(block.getTerminator())) + getOrCreateCondPlaceholder(&block, builder); + } +} + +/// Resolves all SourceOp condition placeholders into NotIOp pass-throughs +/// connected to the real handshake condition values from ShadowCFG. +void ftd::resolveCondPlaceholders(handshake::FuncOp funcOp, OpBuilder &builder, + ftd::ShadowCFG &shadow) { + Block &entryBlock = funcOp.getBody().front(); + + // Collect the condition placeholders (the tagged constants) in the entry. + SmallVector placeholders; + for (Operation &op : entryBlock) { + if (isa(&op) && op.hasAttr(FTD_COND_VAR)) + placeholders.push_back(&op); + } + + for (Operation *ph : placeholders) { + auto bbAttr = ph->getAttrOfType("handshake.bb"); + if (!bbAttr) + continue; + unsigned bbIdx = bbAttr.getUInt(); + + // The real handshake condition this placeholder stands for. + Value realCond = shadow.getCondition(bbIdx); + if (!realCond) + continue; + + if (auto *defOp = realCond.getDefiningOp()) + builder.setInsertionPointAfter(defOp); + else + builder.setInsertionPointToStart(&entryBlock); + + Location loc = ph->getLoc(); + Type chanI1 = ftd::channelifyType(builder.getI1Type()); + + // Forward the real condition through a NotIOp that keeps the placeholder + // tag, so finalizeCondPlaceholders can later short-circuit and remove it. + auto notOp = builder.create(loc, chanI1, realCond); + notOp->setAttr(FTD_COND_VAR, builder.getUnitAttr()); + notOp->setAttr("handshake.bb", bbAttr); + + // Kill the old ConstantOp placeholder and its SourceOp + Operation *phSourceOp = ph->getOperand(0).getDefiningOp(); + ph->getResult(0).replaceAllUsesWith(notOp.getResult()); + ph->erase(); + if (phSourceOp && phSourceOp->use_empty()) + phSourceOp->erase(); } - if (!muxConditionBlock) - muxConditionBlock = muxCondition.getParentBlock(); - return muxConditionBlock; } +/// Short-circuits all NotIOp condition placeholders and erases them. +void ftd::finalizeCondPlaceholders(handshake::FuncOp funcOp) { + SmallVector notOps; + for (auto notOp : funcOp.getOps()) { + if (notOp->hasAttr(FTD_COND_VAR)) + notOps.push_back(notOp); + } + for (auto notOp : notOps) { + notOp.getResult().replaceAllUsesWith(notOp.getOperand()); + + notOp->erase(); + } +} + +// ===--------------------------------------------------------------------=== // +// Static CFG helpers (used only by phi network functions) +// ===--------------------------------------------------------------------=== // + /// Given a block, get its immediate dominator if exists static Block *getImmediateDominator(Region ®ion, Block *bb) { @@ -90,7 +113,7 @@ static Block *getImmediateDominator(Region ®ion, Block *bb) { if (region.getBlocks().empty()) return nullptr; - // The first block in the CFG has both non predecessors and no dominators + // The first block in the CFG has both no predecessors and no dominators if (bb->hasNoPredecessors()) return nullptr; @@ -109,7 +132,7 @@ getDominanceFrontier(Region ®ion) { DenseMap> result; - // Create an empty set of reach available block + // Create an empty dominance-frontier set for each block for (Block &bb : region.getBlocks()) result.insert({&bb, DenseSet()}); @@ -128,7 +151,8 @@ getDominanceFrontier(Region ®ion) { if (numberOfPredecessors < 2) continue; - // Run the algorithm as explained in the paper + // For each predecessor, walk up the dominator tree, adding bb to the + // frontier of every block until bb's immediate dominator is reached. for (auto *pred : predecessors) { Block *runner = pred; // Runner performs a bottom up traversal of the dominator tree @@ -144,148 +168,27 @@ getDominanceFrontier(Region ®ion) { /// Get a list of all the loops in which the consumer is but the producer is /// not, starting from the innermost. -static SmallVector getLoopsConsNotInProd(Block *cons, Block *prod, - mlir::CFGLoopInfo &li) { +static SmallVector getLoopsConsNotInProd(Block *consBlock, + Block *prodBlock, + CFGLoopInfo &loopInfo) { + SmallVector result; - // Get all the loops in which the consumer is but the producer is - // not, starting from the innermost - for (CFGLoop *loop = li.getLoopFor(cons); loop; - loop = loop->getParentLoop()) { - if (!loop->contains(prod)) + CFGLoop *consLoop = loopInfo.getLoopFor(consBlock); + if (!consLoop) + return result; + + // Walk outward from the consumer's innermost loop. + // Collect every loop that does NOT contain the producer. + for (CFGLoop *loop = consLoop; loop; loop = loop->getParentLoop()) { + if (!loop->contains(prodBlock)) result.push_back(loop); } - // Reverse to the get the loops from outermost to innermost + // Reverse to ensure the loops from outermost to innermost std::reverse(result.begin(), result.end()); - return result; -} - -/// Given two sets containing object of type `Block*`, remove the common -/// entries. -static void eliminateCommonBlocks(DenseSet &s1, - DenseSet &s2) { - - SmallVector intersection; - for (auto &e1 : s1) { - if (s2.contains(e1)) - intersection.push_back(e1); - } - - for (auto &bb : intersection) { - s1.erase(bb); - s2.erase(bb); - } -} - -/// Given an operation, returns true if the operation is a conditional branch -/// which terminates a for loop. This is the case if it is in one of the exiting -/// blocks of the innermost loop it is in. -static bool isBranchLoopExit(Operation *op, CFGLoopInfo &li) { - if (isa(op)) { - if (CFGLoop *loop = li.getLoopFor(op->getBlock()); loop) { - llvm::SmallVector exitBlocks; - loop->getExitingBlocks(exitBlocks); - return llvm::find(exitBlocks, op->getBlock()) != exitBlocks.end(); - } - } - return false; -} - -/// Given an operation, return true if the two operands of a multiplexer come -/// from two different loops. When this happens, the mux is connecting two -/// loops. -static bool isaMuxLoop(Operation *mux, CFGLoopInfo &li) { - - auto muxOp = llvm::dyn_cast(mux); - if (!muxOp) - return false; - - auto dataOperands = muxOp.getDataOperands(); - - // Get the basic block of the "real" value, so going up the hierarchy as long - // as there are conditional branches involved. - auto getBasicBlockProducer = [&](Value op) -> Block * { - Block *bb = op.getParentBlock(); - - // If the operand is produced by a real operation, such operation might be a - // conditional branch in the same bb of the original. - if (auto *owner = op.getDefiningOp(); owner) { - while (llvm::isa_and_nonnull(owner) && - owner->getBlock() == muxOp->getBlock()) { - auto op = dyn_cast(owner); - if (op.getOperand(1).getDefiningOp()) { - owner = op.getOperand(1).getDefiningOp(); - bb = owner->getBlock(); - continue; - } - break; - } - } - - return bb; - }; - - return li.getLoopFor(getBasicBlockProducer(dataOperands[0])) != - li.getLoopFor(getBasicBlockProducer(dataOperands[1])); -} - -/// The boolean condition to either generate or suppress a token are computed -/// by considering all the paths from the producer (`start`) to the consumer -/// (`end`). "Each path identifies a Boolean product of elementary conditions -/// expressing the reaching of the target BB from the corresponding member of -/// the set; the product of all such paths are added". -static BoolExpression *enumeratePaths(Block *start, Block *end, - const ftd::BlockIndexing &bi, - const DenseSet &controlDeps) { - // Start with a boolean expression of zero (so that new conditions can be - // added) - BoolExpression *sop = BoolExpression::boolZero(); - - // Find all the paths from the producer to the consumer, using a DFS - std::vector> allPaths = findAllPaths(start, end, bi); - - // If the start and end block are the same (e.g., BB0 to BB0) and there is no - // real path between them, then consider the sop = 1 - if (start == end && allPaths.size() == 0) - sop = BoolExpression::boolOne(); - - // For each path - for (const std::vector &path : allPaths) { - - DenseSet tempCofactorSet; - // Compute the product of the conditions which allow that path to be - // executed - BoolExpression *minterm = - getPathExpression(path, tempCofactorSet, bi, controlDeps, false); - - // Add the value to the result - sop = BoolExpression::boolOr(sop, minterm); - } - return sop->boolMinimizeSop(); -} - -/// Get a boolean expression representing the exit condition of the current -/// loop block. -static BoolExpression *getBlockLoopExitCondition(Block *loopExit, CFGLoop *loop, - CFGLoopInfo &li, - const ftd::BlockIndexing &bi) { - - // Get the boolean expression associated to the block exit - BoolExpression *blockCond = - BoolExpression::parseSop(bi.getBlockCondition(loopExit)); - - // Since we are in a loop, the terminator is a conditional branch. - auto *terminatorOperation = loopExit->getTerminator(); - auto condBranch = dyn_cast(terminatorOperation); - assert(condBranch && "Terminator of a loop must be `cf::CondBranchOp`"); - - // If the destination of the false outcome is not the block, then the - // condition must be negated - if (li.getLoopFor(condBranch.getFalseDest()) != loop) - blockCond->boolNegate(); - return blockCond; + return result; } /// Run the Cytron algorithm to determine, give a set of values, in which blocks @@ -338,6 +241,10 @@ runCrytonAlgorithm(Region &funcRegion, DenseMap &inputBlocks) { return result; } +// ===--------------------------------------------------------------------=== // +// Phi Network +// ===--------------------------------------------------------------------=== // + LogicalResult experimental::ftd::createPhiNetwork( Region &funcRegion, PatternRewriter &rewriter, SmallVector &vals, SmallVector &toSubstitue) { @@ -557,132 +464,20 @@ LogicalResult ftd::createPhiNetworkDeps( return success(); } -/// Starting from a boolean expression which is a single variable (either -/// direct or complement) return its corresponding circuit equivalent. This -/// means, either we obtain the output of the operation determining the -/// condition, or we add a `not` to complement. -static Value boolVariableToCircuit(PatternRewriter &rewriter, - experimental::boolean::BoolExpression *expr, - Block *block, const ftd::BlockIndexing &bi, - bool needsChannelify = true) { - - // Convert the expression into a single condition (for instance, `c0` or - // `~c0`). - SingleCond *singleCond = static_cast(expr); - - // Use the BlockIndexing to access the block corresponding to such condition - // and access its terminator to determine the condition. - auto conditionOpt = bi.getBlockFromCondition(singleCond->id); - if (!conditionOpt.has_value()) - return nullptr; - - auto condition = conditionOpt.value()->getTerminator()->getOperand(0); - - // Add a not if the condition is negated. - if (singleCond->isNegated) { - rewriter.setInsertionPointToStart(block); - auto notIOp = rewriter.create( - block->getOperations().front().getLoc(), - ftd::channelifyType(condition.getType()), condition); - notIOp->setAttr(FTD_OP_TO_SKIP, rewriter.getUnitAttr()); - return notIOp->getResult(0); - } - if (needsChannelify) { - condition.setType(ftd::channelifyType(condition.getType())); - } - return condition; -} - -/// Get a circuit out a boolean expression, depending on the different kinds -/// of expressions you might have. -static Value boolExpressionToCircuit(PatternRewriter &rewriter, - BoolExpression *expr, Block *block, - const ftd::BlockIndexing &bi, - bool needsChannelify = true) { - - // Variable case - if (expr->type == ExpressionType::Variable) - return boolVariableToCircuit(rewriter, expr, block, bi, needsChannelify); - - // Constant case (either 0 or 1) - rewriter.setInsertionPointToStart(block); - auto sourceOp = rewriter.create( - block->getOperations().front().getLoc()); - Value cnstTrigger = sourceOp.getResult(); - - auto intType = rewriter.getIntegerType(1); - auto cstAttr = rewriter.getIntegerAttr( - intType, (expr->type == ExpressionType::One ? 1 : 0)); - - auto constOp = rewriter.create( - block->getOperations().front().getLoc(), cstAttr, cnstTrigger); - - constOp->setAttr(FTD_OP_TO_SKIP, rewriter.getUnitAttr()); - - return constOp.getResult(); -} +// ===--------------------------------------------------------------------=== // +// Regeneration +// ===--------------------------------------------------------------------=== // -/// Convert a `BDD` object as obtained from the bdd expansion to a -/// circuit -static Value bddToCircuit(PatternRewriter &rewriter, BDD *bdd, Block *block, - const ftd::BlockIndexing &bi, - bool needsChannelify = true) { - if (!bdd->successors.has_value()) - return boolExpressionToCircuit(rewriter, bdd->boolVariable, block, bi, - needsChannelify); - - rewriter.setInsertionPointToStart(block); - - // Get the two operands by recursively calling `bddToCircuit` (it possibly - // creates other muxes in a hierarchical way) - SmallVector muxOperands; - muxOperands.push_back(bddToCircuit(rewriter, bdd->successors.value().first, - block, bi, needsChannelify)); - muxOperands.push_back(bddToCircuit(rewriter, bdd->successors.value().second, - block, bi, needsChannelify)); - Value muxCond = boolExpressionToCircuit(rewriter, bdd->boolVariable, block, - bi, needsChannelify); - - // Create the multiplxer and add it to the rest of the circuit - auto muxOp = rewriter.create( - block->getOperations().front().getLoc(), muxOperands[0].getType(), - muxCond, muxOperands); - muxOp->setAttr(FTD_OP_TO_SKIP, rewriter.getUnitAttr()); - - return muxOp.getResult(); -} - -static BoolExpression * -getLoopExitCondition(CFGLoop *loop, std::vector *cofactorList, - mlir::CFGLoopInfo &li, const ftd::BlockIndexing &bi) { - - SmallVector exitBlocks; - loop->getExitingBlocks(exitBlocks); - - BoolExpression *fLoopExit = BoolExpression::boolZero(); - - // Get the list of all the cofactors related to possible exit conditions - for (Block *exitBlock : exitBlocks) { - BoolExpression *blockCond = - getBlockLoopExitCondition(exitBlock, loop, li, bi); - fLoopExit = BoolExpression::boolOr(fLoopExit, blockCond); - cofactorList->push_back(bi.getBlockCondition(exitBlock)); - fLoopExit = fLoopExit->boolMinimize(); - } - - // Sort the cofactors alphabetically - std::sort(cofactorList->begin(), cofactorList->end()); - - return fLoopExit; -} - -void ftd::addRegenOperandConsumer(PatternRewriter &rewriter, +void ftd::addRegenOperandConsumer(mlir::OpBuilder &builder, handshake::FuncOp &funcOp, - Operation *consumerOp, Value operand) { - - mlir::DominanceInfo domInfo; - mlir::CFGLoopInfo loopInfo(domInfo.getDomTree(&funcOp.getBody())); - BlockIndexing bi(funcOp.getBody()); + Operation *consumerOp, Value operand, + ftd::ShadowCFG &shadow) { + + // All analysis runs on the shadow Region (multi-block, real CF terminators) + Region &shadowRegion = shadow.getRegion(); + BlockIndexing bi(shadowRegion); + DominanceInfo domInfo(shadow.shadowFunc); + CFGLoopInfo loopInfo(domInfo.getDomTree(&shadowRegion)); auto startValue = (Value)funcOp.getArguments().back(); // Skip if the consumer was added by this function, if it is an init merge, if @@ -704,6 +499,17 @@ void ftd::addRegenOperandConsumer(PatternRewriter &rewriter, mlir::Operation *producerOp = operand.getDefiningOp(); + uint32_t prodId = 0, consId = 0; + if (operand.getDefiningOp()) + if (auto intAttr = + producerOp->getAttrOfType("handshake.bb")) { + prodId = intAttr.getUInt(); + } + if (auto intAttr = + consumerOp->getAttrOfType("handshake.bb")) { + consId = intAttr.getUInt(); + } + // Skip if the producer was added by this function or if it is an op to skip if (producerOp && (producerOp->hasAttr(FTD_REGEN) || producerOp->hasAttr(FTD_OP_TO_SKIP))) @@ -714,60 +520,66 @@ void ftd::addRegenOperandConsumer(PatternRewriter &rewriter, llvm::isa_and_nonnull(operand.getType())) return; - // Last regenerated value - Value regeneratedValue = operand; + // Map BB indices to shadow blocks for loop analysis + Block *prodBlock = shadow.getBlock(prodId); + Block *consBlock = shadow.getBlock(consId); // Get all the loops for which we need to regenerate the // corresponding value - SmallVector loops = getLoopsConsNotInProd( - consumerOp->getBlock(), operand.getParentBlock(), loopInfo); + SmallVector loops = + getLoopsConsNotInProd(consBlock, prodBlock, loopInfo); unsigned numberOfLoops = loops.size(); - auto cstType = rewriter.getIntegerType(1); + if (numberOfLoops == 0) + return; + + Value regeneratedValue = operand; + auto cstType = builder.getIntegerType(1); auto cstAttr = IntegerAttr::get(cstType, 0); + // The real (flattened) block where new ops are inserted + Block *realBlock = &funcOp.getBody().front(); + auto createRegenMux = [&](CFGLoop *loop) -> handshake::MuxOp { - rewriter.setInsertionPointToStart(loop->getHeader()); - regeneratedValue.setType(channelifyType(regeneratedValue.getType())); + builder.setInsertionPointToStart(realBlock); + + // BB index of the loop header (for handshake.bb tagging) + unsigned headerBBIdx = shadow.getBlockIndex(loop->getHeader()); + auto headerBBAttr = ftd::getBBIndexAttr(builder.getContext(), headerBBIdx); // Determine the loop exit condition: - // - If the condition spans multiple cofactors, build a BDD and - // translate it into a circuit. - // - Otherwise, use the simple terminating condition of the exiting block Value conditionValue; - std::vector cofactorList; - BoolExpression *exitCondition = - getLoopExitCondition(loop, &cofactorList, loopInfo, bi); - if (size(cofactorList) > 1) { - BDD *bdd = buildBDD(exitCondition, cofactorList); - conditionValue = - bddToCircuit(rewriter, bdd, loop->getHeader(), bi, false); - } else - conditionValue = loop->getExitingBlock()->getTerminator()->getOperand(0); + Block *loopHeader = loop->getHeader(); + + conditionValue = computeLoopBackedgeCondition( + builder, loopHeader, realBlock, bi, nullptr, &shadow); // Create the false constant to feed `init` - auto constOp = rewriter.create(consumerOp->getLoc(), - cstAttr, startValue); - constOp->setAttr(FTD_INIT_MERGE, rewriter.getUnitAttr()); + auto constOp = builder.create(consumerOp->getLoc(), + cstAttr, startValue); + constOp->setAttr(FTD_INIT_MERGE, builder.getUnitAttr()); + constOp->setAttr("handshake.bb", headerBBAttr); + + Operation *initOp; + initOp = + builder.create(consumerOp->getLoc(), conditionValue); - // Create the `init` operation - SmallVector mergeOperands = {constOp.getResult(), conditionValue}; - auto initMergeOp = rewriter.create(consumerOp->getLoc(), - mergeOperands); - initMergeOp->setAttr(FTD_INIT_MERGE, rewriter.getUnitAttr()); + initOp->setAttr(FTD_INIT_MERGE, builder.getUnitAttr()); + initOp->setAttr("handshake.bb", headerBBAttr); // The multiplexer is to be fed by the init block, and takes as inputs the // regenerated value and the result itself (to be set after) it was created. - auto selectSignal = initMergeOp.getResult(); + auto selectSignal = initOp->getResult(0); selectSignal.setType(channelifyType(selectSignal.getType())); SmallVector muxOperands = {regeneratedValue, regeneratedValue}; - auto muxOp = rewriter.create(regeneratedValue.getLoc(), - regeneratedValue.getType(), - selectSignal, muxOperands); + auto muxOp = builder.create(regeneratedValue.getLoc(), + regeneratedValue.getType(), + selectSignal, muxOperands); muxOp->setOperand(2, muxOp->getResult(0)); - muxOp->setAttr(FTD_REGEN, rewriter.getUnitAttr()); + muxOp->setAttr(FTD_REGEN, builder.getUnitAttr()); + muxOp->setAttr("handshake.bb", headerBBAttr); return muxOp; }; @@ -789,227 +601,16 @@ void ftd::addRegenOperandConsumer(PatternRewriter &rewriter, consumerOp->replaceUsesOfWith(operand, regeneratedValue); } -// Returns true if loop is a while loop, detected by the loop header being -// also a loop exit and not a loop latch -static bool isWhileLoop(CFGLoop *loop) { - if (!loop) - return false; - - Block *headerBlock = loop->getHeader(); - - SmallVector exitBlocks; - loop->getExitingBlocks(exitBlocks); - - SmallVector latchBlocks; - loop->getLoopLatches(latchBlocks); - - return llvm::is_contained(exitBlocks, headerBlock) && - !llvm::is_contained(latchBlocks, headerBlock); -} - -using PairOperandConsumer = std::pair; - -// Find the closest loop exit between a producer inside the loop and a consumer -// outside the loop. -static Block *findClosestLoopExit(Operation *consumer, Value connection, - const ftd::BlockIndexing &bi, - SmallVector exitBlocks) { - // Find all the paths from the producer to the consumer using DFS - std::vector> allPaths = - findAllPaths(connection.getParentBlock(), consumer->getBlock(), bi); - - Block *closestExit = nullptr; - unsigned minDistance = std::numeric_limits::max(); - - // For every path from producer to consumer , check if it passes through any - // of the loop’s exit blocks. - for (const auto &path : allPaths) { - bool foundInThisPath = false; - for (unsigned i = 0; i < path.size() && !foundInThisPath; ++i) { - Block *pathBlock = path[i]; - for (Block *exitBlock : exitBlocks) { - if (pathBlock == exitBlock) { - // Update the closest exit if this one is nearer to the producer - if (i < minDistance) { - minDistance = i; - closestExit = exitBlock; - } - foundInThisPath = true; - break; - } - } - } - } - - assert(closestExit && - "No loop exit found in any path between producer and consumer."); - return closestExit; -} - -/// Insert a branch to the correct position, taking into account whether it -/// should work to suppress the over-production of tokens or self-regeneration -static Value addSuppressionInLoop(PatternRewriter &rewriter, CFGLoop *loop, - Operation *consumer, Value connection, - BranchToLoopType btlt, CFGLoopInfo &li, - std::vector &toCover, - const ftd::BlockIndexing &bi) { - - handshake::ConditionalBranchOp branchOp; - - // Do not add the branch in case of a while loop with backward edge - if (btlt == BackwardRelationship && isWhileLoop(loop)) - return connection; - - std::vector cofactorList; - SmallVector exitBlocks; - loop->getExitingBlocks(exitBlocks); - BoolExpression *fLoopExit = getLoopExitCondition(loop, &cofactorList, li, bi); - // Choose the closest loop exit to the producer and place the suppression - // there - Block *loopExit = findClosestLoopExit(consumer, connection, bi, exitBlocks); - - // Apply a BDD expansion to the loop exit expression and the list of - // cofactors - BDD *bdd = buildBDD(fLoopExit, cofactorList); - - // Convert the boolean expression obtained through BDD to a circuit - Value branchCond = bddToCircuit(rewriter, bdd, loopExit, bi); - - Operation *loopTerminator = loopExit->getTerminator(); - assert(isa(loopTerminator) && - "Terminator condition of a loop exit must be a conditional " - "branch."); - - rewriter.setInsertionPointToStart(loopExit); - - branchOp = rewriter.create( - loopExit->getOperations().front().getLoc(), - ftd::getListTypes(connection.getType()), branchCond, connection); - - Value newConnection = btlt == MoreProducerThanConsumers - ? branchOp.getTrueResult() - : branchOp.getFalseResult(); - - // If we are handling a case with more producers than consumers, the new - // branch must undergo the `addSupp` function so we add it to our structure - // to be able to loop over it - if (btlt == MoreProducerThanConsumers) { - branchOp->setAttr(FTD_NEW_SUPP, rewriter.getUnitAttr()); - toCover.emplace_back(newConnection, consumer); - } +// ===--------------------------------------------------------------------=== // +// Suppression dispatch +// ===--------------------------------------------------------------------=== // - consumer->replaceUsesOfWith(connection, newConnection); - return newConnection; -} - -/// Apply the algorithm from FPL'22 to handle a non-loop situation of -/// producer and consumer -static void insertDirectSuppression( - PatternRewriter &rewriter, handshake::FuncOp &funcOp, Operation *consumer, - Value connection, const ftd::BlockIndexing &bi, - ControlDependenceAnalysis::BlockControlDepsMap &cdAnalysis) { - - Block *entryBlock = &funcOp.getBody().front(); - Block *producerBlock = connection.getParentBlock(); - Block *consumerBlock = consumer->getBlock(); - Value muxCondition = nullptr; - - // Account for the condition of a Mux only if it corresponds to a GAMMA GSA - // gate and the producer is one of its data inputs - bool accountMuxCondition = llvm::isa(consumer) && - consumer->hasAttr(FTD_EXPLICIT_GAMMA) && - (consumer->getOperand(1) == connection || - consumer->getOperand(2) == connection); - - // Get the control dependencies from the producer - DenseSet prodControlDeps = - cdAnalysis[producerBlock].forwardControlDeps; - - // Get the control dependencies from the consumer - DenseSet consControlDeps = - cdAnalysis[consumer->getBlock()].forwardControlDeps; - - // If the mux condition is to be taken into account, then the control - // dependencies of the mux conditions are to be added to the consumer control - // dependencies - if (accountMuxCondition) { - muxCondition = consumer->getOperand(0); - Block *muxConditionBlock = returnMuxConditionBlock(muxCondition); - DenseSet condControlDeps = - cdAnalysis[muxConditionBlock].forwardControlDeps; - for (auto &x : condControlDeps) - consControlDeps.insert(x); - } - - // Get rid of common entries in the two sets - eliminateCommonBlocks(prodControlDeps, consControlDeps); - - // Compute the activation function of producer and consumer - BoolExpression *fProd = - enumeratePaths(entryBlock, producerBlock, bi, prodControlDeps); - BoolExpression *fCons = - enumeratePaths(entryBlock, consumerBlock, bi, consControlDeps); - - if (accountMuxCondition) { - Block *muxConditionBlock = returnMuxConditionBlock(muxCondition); - BoolExpression *selectOperandCondition = - BoolExpression::parseSop(bi.getBlockCondition(muxConditionBlock)); - - // The condition must be taken into account for `fCons` only if the - // producer is not control dependent from the block which produces the - // condition of the mux - if (!prodControlDeps.contains(muxConditionBlock)) { - if (consumer->getOperand(1) == connection) - fCons = BoolExpression::boolAnd(fCons, - selectOperandCondition->boolNegate()); - else - fCons = BoolExpression::boolAnd(fCons, selectOperandCondition); - } - } - - /// f_supp = f_prod and not f_cons - BoolExpression *fSup = BoolExpression::boolAnd(fProd, fCons->boolNegate()); - fSup = fSup->boolMinimize(); - - // If the activation function is not zero, then a suppress block is to be - // inserted - if (fSup->type != experimental::boolean::ExpressionType::Zero) { - std::set blocks = fSup->getVariables(); - - std::vector cofactorList(blocks.begin(), blocks.end()); - BDD *bdd = buildBDD(fSup, cofactorList); - Value branchCond = bddToCircuit(rewriter, bdd, consumer->getBlock(), bi); - - rewriter.setInsertionPointToStart(consumer->getBlock()); - auto branchOp = rewriter.create( - consumer->getLoc(), ftd::getListTypes(connection.getType()), branchCond, - connection); - - // Take into account the possibility of a mux to get the condition input - // also as data input. In this case, a branch needs to be created, but only - // the corresponding data input is affected. The conditions below take into - // account this possibility. - for (auto &use : connection.getUses()) { - if (use.getOwner() != consumer) - continue; - if (llvm::isa(consumer) && use.getOperandNumber() == 0) - continue; - use.set(branchOp.getFalseResult()); - } - } -} - -void ftd::addSuppOperandConsumer(PatternRewriter &rewriter, +void ftd::addSuppOperandConsumer(mlir::OpBuilder &builder, handshake::FuncOp &funcOp, - Operation *consumerOp, Value operand) { + Operation *consumerOp, Value operand, + ShadowCFG &shadow) { - Region ®ion = funcOp.getBody(); - mlir::DominanceInfo domInfo; - mlir::CFGLoopInfo loopInfo(domInfo.getDomTree(®ion)); - BlockIndexing bi(region); - auto cda = ControlDependenceAnalysis(region).getAllBlockDeps(); - - // Skip the prod-cons if the producer is part of the operations related to + // Skip the prod-cons if the consumer is part of the operations related to // the BDD expansion or INIT merges if (consumerOp->hasAttr(FTD_OP_TO_SKIP) || consumerOp->hasAttr(FTD_INIT_MERGE)) @@ -1020,37 +621,41 @@ void ftd::addSuppOperandConsumer(PatternRewriter &rewriter, consumerOp->getOperand(0) != operand) return; - // The consumer block is the block which contains the consumer - Block *consumerBlock = consumerOp->getBlock(); + // Read BB indices from handshake.bb attributes + unsigned consBBIdx = 0; + if (auto attr = consumerOp->getAttrOfType("handshake.bb")) + consBBIdx = attr.getUInt(); + + unsigned prodBBIdx = 0; + if (Operation *producerOp = operand.getDefiningOp()) + if (auto attr = producerOp->getAttrOfType("handshake.bb")) + prodBBIdx = attr.getUInt(); - // The producer block is the block which contains the producer, and it - // corresponds to the parent block of the operand. Since the operand might - // have no producer operation (if it is a function argument) then this is the - // only way to get the relevant information. - Block *producerBlock = operand.getParentBlock(); + // Map to shadow blocks for analysis + Block *consumerBlock = shadow.getBlock(consBBIdx); + Block *producerBlock = shadow.getBlock(prodBBIdx); // If the consumer and the producer are in the same block without the // consumer being a multiplexer skip because no delivery is needed if (consumerBlock == producerBlock && - !llvm::isa(consumerOp)) + (!llvm::isa(consumerOp) || + operand.getDefiningOp()->hasAttr(FTD_EXPLICIT_GAMMA))) { return; + } if (Operation *producerOp = operand.getDefiningOp(); producerOp) { - // A conditional branch should undergo the suppression mechanism only if it - // has the `FTD_NEW_SUPP` annotation, set in `addMoreSuppressionInLoop`. In - // any other cases, suppressing a branch ends up with incorrect results. - if (llvm::isa(producerOp) && - !producerOp->hasAttr(FTD_NEW_SUPP)) + // In any cases, suppressing a branch ends up with incorrect results. + if (llvm::isa(producerOp)) return; - // Skip the prod-cons if the consumer is part of the operations + // Skip the prod-cons if the producer is part of the operations // related to the BDD expansion or INIT merges if (producerOp->hasAttr(FTD_OP_TO_SKIP) || producerOp->hasAttr(FTD_INIT_MERGE)) return; - // Skip if either the producer of the consumer are + // Skip if either the producer or the consumer are // related to memory operations, or if the consumer is a conditional // branch if (llvm::isa_and_nonnull(consumerOp) || @@ -1060,7 +665,6 @@ void ftd::addSuppOperandConsumer(PatternRewriter &rewriter, llvm::isa_and_nonnull(producerOp) || llvm::isa_and_nonnull(consumerOp) || llvm::isa_and_nonnull(consumerOp) || - llvm::isa_and_nonnull(consumerOp) || llvm::isa_and_nonnull(consumerOp) || (llvm::isa(consumerOp) && !llvm::isa(consumerOp)) || @@ -1069,74 +673,21 @@ void ftd::addSuppOperandConsumer(PatternRewriter &rewriter, llvm::isa(operand.getType())) return; - // The next step is to identify the relationship between the producer - // and consumer in hand: Are they in the same loop or at different - // loop levels? Are they connected through a backward edge? - - // Set true if the producer is in a loop which does not contains - // the consumer - bool producingGtUsing = - loopInfo.getLoopFor(producerBlock) && - !loopInfo.getLoopFor(producerBlock)->contains(consumerBlock); - - auto *consumerLoop = loopInfo.getLoopFor(consumerBlock); - std::vector newToCover; - - // Set to true if the consumer uses its own result - bool selfRegeneration = - llvm::any_of(consumerOp->getResults(), - [&operand](const Value &v) { return v == operand; }); - - // We need to suppress all the tokens produced within a loop and - // used outside each time the loop is not terminated. This should be - // done for as many loops there are - if (producingGtUsing && !isBranchLoopExit(producerOp, loopInfo)) { - Value con = operand; - for (CFGLoop *loop = loopInfo.getLoopFor(producerBlock); loop; - loop = loop->getParentLoop()) { - - // For each loop containing the producer but not the consumer, add - // the branch - if (!loop->contains(consumerBlock)) - con = addSuppressionInLoop(rewriter, loop, consumerOp, con, - MoreProducerThanConsumers, loopInfo, - newToCover, bi); - } - - for (auto &pair : newToCover) - addSuppOperandConsumer(rewriter, funcOp, pair.second, pair.first); - - return; - } - - // We need to suppress a token if the consumer is the producer itself - // within a loop - if (selfRegeneration && consumerLoop && - !producerOp->hasAttr(FTD_NEW_SUPP)) { - addSuppressionInLoop(rewriter, consumerLoop, consumerOp, operand, - SelfRegeneration, loopInfo, newToCover, bi); + // Skip cf::CondBranchOp consumers unless this operand is the condition + // input (operand 0) of the block's terminator. + if (llvm::isa_and_nonnull(consumerOp) && + (consumerOp != consumerBlock->getTerminator() || + operand != consumerOp->getOperand(0))) return; - } - // We need to suppress a token if the consumer comes before the - // producer (backward edge) - if ((bi.isGreater(producerBlock, consumerBlock) || - (llvm::isa(consumerOp) && - producerBlock == consumerBlock && - isaMuxLoop(consumerOp, loopInfo))) && - consumerLoop) { - addSuppressionInLoop(rewriter, consumerLoop, consumerOp, operand, - BackwardRelationship, loopInfo, newToCover, bi); - return; - } + // Handle the suppression in all the other cases (including the operand + // being a function argument) + insertDirectSuppression(builder, funcOp, consumerOp, operand, shadow); } - - // Handle the suppression in all the other cases (including the operand being - // a function argument) - insertDirectSuppression(rewriter, funcOp, consumerOp, operand, bi, cda); } -void ftd::addSupp(handshake::FuncOp &funcOp, PatternRewriter &rewriter) { +void ftd::addSupp(handshake::FuncOp &funcOp, mlir::OpBuilder &builder, + ShadowCFG &shadow) { // Set of original operations in the IR std::vector consumersToCover; @@ -1145,11 +696,12 @@ void ftd::addSupp(handshake::FuncOp &funcOp, PatternRewriter &rewriter) { for (auto *consumerOp : consumersToCover) { for (auto operand : consumerOp->getOperands()) - addSuppOperandConsumer(rewriter, funcOp, consumerOp, operand); + addSuppOperandConsumer(builder, funcOp, consumerOp, operand, shadow); } } -void ftd::addRegen(handshake::FuncOp &funcOp, PatternRewriter &rewriter) { +void ftd::addRegen(handshake::FuncOp &funcOp, mlir::OpBuilder &builder, + ShadowCFG &shadow) { // Set of original operations in the IR std::vector consumersToCover; @@ -1159,15 +711,19 @@ void ftd::addRegen(handshake::FuncOp &funcOp, PatternRewriter &rewriter) { // For each producer/consumer relationship for (Operation *consumerOp : consumersToCover) { for (Value operand : consumerOp->getOperands()) - addRegenOperandConsumer(rewriter, funcOp, consumerOp, operand); + addRegenOperandConsumer(builder, funcOp, consumerOp, operand, shadow); } } -LogicalResult experimental::ftd::addGsaGates(Region ®ion, - PatternRewriter &rewriter, - const gsa::GSAAnalysis &gsa, - Backedge startValue, - bool removeTerminators) { +// ===--------------------------------------------------------------------=== // +// GSA Gates +// ===--------------------------------------------------------------------=== // + +LogicalResult experimental::ftd::addGsaGates( + Region ®ion, PatternRewriter &rewriter, const gsa::GSAAnalysis &gsa, + Backedge startValue, + DenseMap> *pendingMuxOperands, + bool removeTerminators) { using namespace experimental::gsa; BlockIndexing bi(region); @@ -1185,6 +741,25 @@ LogicalResult experimental::ftd::addGsaGates(Region ®ion, // To simplify the way GSA functions are handled, each of them has an unique // index. + // This function operates in two modes: + // (1) Called from handshake transformation passes where all values are + // already handshake channels. + // (2) Called early in CF-to-handshake conversion + // when some CF values are still present. + // The last two parameters distinguish the modes: for (1) they are + // nullptr/false, for (2) they are non-null/true. + // + // In mode (1), no special management is needed: Muxes are inserted with all + // connections finalized immediately. + // + // In mode (2), Mux operands are created as backedge placeholders. We maintain + // a side structure mapping these placeholders to their corresponding + // handshake values, which allows us to replace the backedges later, once the + // handshake values are fully finalized. In case of Mux fed from a Mux tree, + // the pendingMuxOperands is propoagated to the Mux decomposition tree to + // store all cf values involved. In case the Mux is fed from a Merge only the + // cf values feeding the Merge will be pushed to the pendingMuxOperands + struct MissingGsa { // Index of the GSA function to modify unsigned phiIndex; @@ -1223,7 +798,7 @@ LogicalResult experimental::ftd::addGsaGates(Region ®ion, // Checks whether one index is empty int nullOperand = -1; - // For each of its operand + // For each of its operands for (auto *operand : gate->operands) { // If the input is another GSA function, then a dummy value is used as // operand and the operations will be reconnected later on. @@ -1239,26 +814,58 @@ LogicalResult experimental::ftd::addGsaGates(Region ®ion, operands.emplace_back(nullptr); } else { auto val = std::get(operand->input); - operands.emplace_back(val); + // If there is no risk that values are not finalized yet, + // pendingMuxOperands will be a nullptr + if (pendingMuxOperands == nullptr) + operands.emplace_back(val); + else { + // Create backedge of the same type + BackedgeBuilder beb(rewriter, block.front().getLoc()); + Backedge be = beb.get(ftd::channelifyType(val.getType())); + // Use backedge as mux operand + operands.emplace_back(be); + // Remember how to resolve it later + (*pendingMuxOperands)[val].push_back(be); + } } operandIndex++; } // Get the condition for the block exiting - // Determine the gate exit condition: - // - If the condition spans multiple cofactors, build a BDD and - // translate it into a circuit. - // - Otherwise, use the simple terminating condition of the block Value conditionValue; - if (size(gate->cofactorList) > 1) { - // Apply a BDD expansion to the loop exit expression and the list of - // cofactors - BDD *bdd = buildBDD(gate->condition, gate->cofactorList); - // Convert the boolean expression obtained through BDD to a circuit - conditionValue = - bddToCircuit(rewriter, bdd, gate->getBlock(), bi, false); - } else - conditionValue = gate->conditionBlock->getTerminator()->getOperand(0); + + // Determine the gate exit condition + if (gate->gsaGateFunction == MuGate) { + // For MU gates, we generate the condition based on the + // reaching condition from the loop header back to itself. + Block *loopHeader = gate->getBlock(); + conditionValue = computeLoopBackedgeCondition( + rewriter, loopHeader, loopHeader, bi, pendingMuxOperands); + + } else { + // [Gamma Logic] + if (size(gate->cofactorList) > 1) { + // Apply a BDD expansion to the loop exit expression and the list of + // cofactors + BDD *bdd = buildBDD(gate->condition, gate->cofactorList); + // Convert the boolean expression obtained through BDD to a circuit + // We pass an empty registry, since this is not an expression for + // suppression and does not require distribution. + SignalRegistry emptyRegistry; + conditionValue = + bddToCircuit(rewriter, bdd, gate->getBlock(), emptyRegistry, {}, + bi, pendingMuxOperands); + } else { + // Use a SourceOp placeholder for the condition value, which will be + // replaced with the actual condition input after every suppression + // is done. + conditionValue = + getOrCreateCondPlaceholder(gate->conditionBlock, rewriter); + // Ensure type consistency (Channel vs i1) + if (!conditionValue.getType().isa()) + conditionValue.setType(channelifyType(conditionValue.getType())); + } + } // If the function is MU, then we create a merge // and use its result as condition @@ -1266,31 +873,15 @@ LogicalResult experimental::ftd::addGsaGates(Region ®ion, mlir::DominanceInfo domInfo; mlir::CFGLoopInfo loopInfo(domInfo.getDomTree(®ion)); - // The inputs of the merge are the condition value and a `false` - // constant driven by the start value of the function. This will - // created later on, so we use a dummy value. - SmallVector mergeOperands; - mergeOperands.push_back(conditionValue); - mergeOperands.push_back(conditionValue); - - auto initMergeOp = - rewriter.create(loc, mergeOperands); + Operation *initOp; + initOp = rewriter.create(loc, conditionValue); - initMergeOp->setAttr(FTD_INIT_MERGE, rewriter.getUnitAttr()); + initOp->setAttr(FTD_INIT_MERGE, rewriter.getUnitAttr()); + setBBAttr(initOp, gate->getBlock(), rewriter); // Replace the new condition value - conditionValue = initMergeOp->getResult(0); + conditionValue = initOp->getResult(0); conditionValue.setType(channelifyType(conditionValue.getType())); - - // Add the activation constant driven by the backedge value, which will - // be then updated with the real start value, once available - auto cstType = rewriter.getIntegerType(1); - auto cstAttr = IntegerAttr::get(cstType, 0); - rewriter.setInsertionPointToStart(initMergeOp->getBlock()); - auto constOp = rewriter.create( - initMergeOp->getLoc(), cstAttr, startValue); - constOp->setAttr(FTD_INIT_MERGE, rewriter.getUnitAttr()); - initMergeOp->setOperand(0, constOp.getResult()); } // When a single input gamma is encountered, a mux is inserted as a @@ -1302,10 +893,11 @@ LogicalResult experimental::ftd::addGsaGates(Region ®ion, } // Create the multiplexer - auto mux = rewriter.create(loc, gate->result.getType(), - conditionValue, operands); + auto mux = rewriter.create( + loc, ftd::channelifyType(gate->result.getType()), conditionValue, + operands); - // The one input gamma is marked at an operation to skip in the IR and + // The one input gamma is marked as an operation to skip in the IR and // later removed if (nullOperand >= 0) oneInputGammaList.insert(mux); @@ -1345,9 +937,112 @@ LogicalResult experimental::ftd::addGsaGates(Region ®ion, ? 1 : 2; op->getResult(0).replaceAllUsesWith(op->getOperand(operandToUse)); + + for (auto &[idx, mappedOp] : gsaList) { + if (mappedOp == op) + mappedOp = nullptr; + } rewriter.eraseOp(op); } + // Simplify the generated GSA Mux tree by applying common subexpression + // elimination and reduction rules from the bottom up. A reverse mapping + // resolves temporary backedge placeholders to their original CFG values + // ensuring equivalent inputs are correctly identified across different + // branches. + DenseMap backedgeToOriginal; + if (pendingMuxOperands) { + for (auto &[orig, bes] : *pendingMuxOperands) { + for (auto &be : bes) { + backedgeToOriginal[Value(be)] = orig; + } + } + } + + // Helper functions evaluate the structural equivalence of two values. + // Equivalence is based on Value identity (SSA identity): two Values are + // equivalent iff they are the same SSA wire. Backedge placeholders are + // resolved to their original CF values first. + auto getEffectiveValue = [&](Value v) -> Value { + if (!v) + return v; + auto it = backedgeToOriginal.find(v); + if (it != backedgeToOriginal.end()) + return it->second; + return v; + }; + + auto areEquivalentValues = [&](Value a, Value b) { + if (a == b) + return true; + if (!a || !b) + return false; + + Value effA = getEffectiveValue(a); + Value effB = getEffectiveValue(b); + return effA == effB; + }; + + // Iterate through the generated operations until the tree structure fully + // converges. The first step performs common subexpression elimination by + // scanning horizontally across all generated multiplexers. Whenever two + // distinct multiplexers share identical selection conditions and data + // inputs, they are merged into a single operation to eliminate structural + // duplication. + bool changed = true; + while (changed) { + changed = false; + + for (auto it1 = gsaList.begin(); it1 != gsaList.end(); ++it1) { + if (!it1->second) + continue; + auto mux1 = dyn_cast(it1->second); + if (!mux1 || mux1.getNumOperands() != 3) + continue; + + for (auto it2 = std::next(it1); it2 != gsaList.end(); ++it2) { + if (!it2->second) + continue; + auto mux2 = dyn_cast(it2->second); + if (!mux2 || mux2.getNumOperands() != 3) + continue; + + if (areEquivalentValues(mux1.getSelectOperand(), + mux2.getSelectOperand()) && + areEquivalentValues(mux1.getDataOperands()[0], + mux2.getDataOperands()[0]) && + areEquivalentValues(mux1.getDataOperands()[1], + mux2.getDataOperands()[1])) { + + mux2.getResult().replaceAllUsesWith(mux1.getResult()); + rewriter.eraseOp(mux2); + it2->second = nullptr; + changed = true; + } + } + } + + // The second step applies the reduction rule by scanning vertically + // through the tree. Any multiplexer whose true and false data inputs + // resolve to the same underlying value is fundamentally redundant and + // is bypassed entirely by routing its data input directly to its users. + for (auto &[idx, op] : gsaList) { + if (!op) + continue; + auto mux = dyn_cast(op); + if (!mux || mux.getNumOperands() != 3) + continue; + + if (areEquivalentValues(mux.getDataOperands()[0], + mux.getDataOperands()[1])) { + mux.getResult().replaceAllUsesWith(mux.getDataOperands()[0]); + rewriter.eraseOp(mux); + op = nullptr; + changed = true; + } + } + } + if (!removeTerminators) return success(); @@ -1394,7 +1089,7 @@ LogicalResult ftd::replaceMergeToGSA(handshake::FuncOp &funcOp, continue; gsa::GSAAnalysis gsa(merge, funcOp.getRegion()); if (failed(ftd::addGsaGates(funcOp.getRegion(), rewriter, gsa, - startValueBackedge, false))) + startValueBackedge, nullptr, false))) return failure(); // Get rid of the merge diff --git a/experimental/lib/Support/FtdSupport.cpp b/experimental/lib/Support/FtdSupport.cpp index f4e85b2593..12eeef9a67 100644 --- a/experimental/lib/Support/FtdSupport.cpp +++ b/experimental/lib/Support/FtdSupport.cpp @@ -6,9 +6,8 @@ // //===----------------------------------------------------------------------===// // -// Implements some utility functions which are useful for both the fast token -// delivery algorithm and for the GSA analysis pass. All the functions are about -// analyzing relationships between blocks and handshake operations. +// Implements utility functions shared across the FTD algorithm: CFG analysis, +// type helpers, IR attribute utilities, and condition placeholders. // //===----------------------------------------------------------------------===// @@ -22,6 +21,10 @@ using namespace mlir; using namespace dynamatic; using namespace dynamatic::experimental; +// ===--------------------------------------------------------------------=== // +// BlockIndexing +// ===--------------------------------------------------------------------=== // + std::string ftd::BlockIndexing::getBlockCondition(Block *block) const { return "c" + std::to_string(getIndexFromBlock(block).value_or(0)); } @@ -34,12 +37,12 @@ ftd::BlockIndexing::BlockIndexing(Region ®ion) { for (Block &bb : region.getBlocks()) allBlocks.push_back(&bb); - // Sort the vector according to the dominance information, so that a block - // comes before each dominators. + // Sort the vector according to the dominance information, so that each + // block comes after its dominators. llvm::sort(allBlocks.begin(), allBlocks.end(), [&](Block *a, Block *b) { return domInfo.dominates(a, b); }); - // Associate a smalled index in the map to the blocks at higer levels of the + // Associate a smaller index in the map to the blocks at higher levels of the // dominance tree for (auto [blockID, bb] : llvm::enumerate(allBlocks)) { indexToBlock.insert({blockID, bb}); @@ -57,6 +60,7 @@ ftd::BlockIndexing::getBlockFromIndex(unsigned index) const { std::optional ftd::BlockIndexing::getBlockFromCondition(StringRef condition) const { + // Conditions are named "cN"; drop the leading 'c' and parse the index N. std::string conditionNumber = condition.str(); conditionNumber.erase(0, 1); StringRef conditionRef = conditionNumber; @@ -90,6 +94,10 @@ bool ftd::BlockIndexing::isLess(Block *bb1, Block *bb2) const { return index1 < index2; } +// ===--------------------------------------------------------------------=== // +// Loop analysis utilities +// ===--------------------------------------------------------------------=== // + /// Recursively check whether 2 blocks belong to the same loop, starting /// from the inner-most loops. static bool isSameLoop(const CFGLoop *loop1, const CFGLoop *loop2) { @@ -105,7 +113,7 @@ bool ftd::isSameLoopBlocks(Block *source, Block *dest, return isSameLoop(li.getLoopFor(source), li.getLoopFor(dest)); } -//// Recursively check whether `inner` is the same loop as `outer`, or is +/// Recursively check whether `inner` is the same loop as `outer`, or is /// nested inside `outer` (i.e., any ancestor of `inner` matches `outer`). static bool isSameOrInnerLoop(const CFGLoop *inner, const CFGLoop *outer) { if (!inner || !outer) @@ -118,6 +126,10 @@ bool ftd::isSameOrInnerLoopBlocks(Block *source, Block *dest, return isSameOrInnerLoop(li.getLoopFor(source), li.getLoopFor(dest)); } +// ===--------------------------------------------------------------------=== // +// Path finding +// ===--------------------------------------------------------------------=== // + /// Recursive function which allows to obtain all the paths from block `start` /// to block `end` using a DFS. static void dfsAllPaths(Block *start, Block *end, std::vector &path, @@ -144,7 +156,7 @@ static void dfsAllPaths(Block *start, Block *end, std::vector &path, // Else, for each successor which was not visited, run DFS again for (Block *successor : start->getSuccessors()) { - // Do not run DFS if the successor is in the list of blocks to traverse + // Do not run DFS into a successor that is in the list of blocks to avoid bool incorrectPath = false; for (auto *toAvoid : blocksToAvoid) { if (toAvoid == successor && bi.isGreater(toAvoid, blockToTraverse)) { @@ -182,6 +194,10 @@ ftd::findAllPaths(Block *start, Block *end, const BlockIndexing &bi, return allPaths; } +// ===--------------------------------------------------------------------=== // +// Path expression generation +// ===--------------------------------------------------------------------=== // + boolean::BoolExpression * ftd::getPathExpression(ArrayRef path, DenseSet &blockIndexSet, @@ -220,7 +236,7 @@ ftd::getPathExpression(ArrayRef path, boolean::BoolExpression *pathCondition = boolean::BoolExpression::parseSop(blockCondition); - // Possibly add the condition to the list of cofactors + // Add the condition's block index to the list of cofactors blockIndexSet.insert(blockIndex); // Negate the condition if `secondBlock` is reached when the condition @@ -235,10 +251,43 @@ ftd::getPathExpression(ArrayRef path, } } - // Minimize the condition and return + // Return the path condition (minimization is left to the caller) return exp; } +// ===--------------------------------------------------------------------=== // +// Reachability +// ===--------------------------------------------------------------------=== // + +bool ftd::isReachable(Block *start, Block *end) { + if (start == end) + return true; + + DenseSet visited; + SmallVector stack; + stack.push_back(start); + visited.insert(start); + + while (!stack.empty()) { + Block *curr = stack.pop_back_val(); + + if (curr == end) + return true; + + for (Block *succ : curr->getSuccessors()) { + if (!visited.count(succ)) { + visited.insert(succ); + stack.push_back(succ); + } + } + } + return false; +} + +// ===--------------------------------------------------------------------=== // +// Type utilities +// ===--------------------------------------------------------------------=== // + Type ftd::channelifyType(Type type) { return llvm::TypeSwitch(type) .Case( @@ -259,3 +308,63 @@ Type ftd::channelifyType(Type type) { SmallVector ftd::getListTypes(Type inputType, unsigned size) { return SmallVector(size, channelifyType(inputType)); } + +// ===--------------------------------------------------------------------=== // +// IR attribute utilities +// ===--------------------------------------------------------------------=== // + +IntegerAttr ftd::getBBIndexAttr(MLIRContext *ctx, unsigned bbIdx) { + return IntegerAttr::get(IntegerType::get(ctx, 32, IntegerType::Unsigned), + bbIdx); +} + +void ftd::setBBAttr(Operation *op, Block *block, OpBuilder &builder) { + unsigned idx = 0; + for (Block &blk : *block->getParent()) { + if (&blk == block) + break; + idx++; + } + op->setAttr("handshake.bb", getBBIndexAttr(builder.getContext(), idx)); +} + +void ftd::setBBAttr(Operation *op, IntegerAttr bbAttr) { + if (bbAttr) + op->setAttr("handshake.bb", bbAttr); +} + +void ftd::setBBAttrWithFallback(Operation *op, IntegerAttr bbAttr, Block *block, + OpBuilder &builder) { + if (bbAttr) + return setBBAttr(op, bbAttr); + setBBAttr(op, block, builder); +} + +// ===--------------------------------------------------------------------=== // +// Condition placeholder +// ===--------------------------------------------------------------------=== // + +Value ftd::getOrCreateCondPlaceholder(Block *condBlock, OpBuilder &builder) { + // Reuse the placeholder already created for this block, if any. + for (Operation &op : *condBlock) { + if (isa(&op) && op.hasAttr(FTD_COND_VAR)) + return op.getResult(0); + } + // Otherwise create a source -> constant pair standing in for the block's + // condition. The constant carries FTD_COND_VAR (it is the placeholder), and + // both ops carry FTD_OP_TO_SKIP so the FTD passes leave them alone. + OpBuilder::InsertionGuard guard(builder); + builder.setInsertionPointToStart(condBlock); + auto sourceOp = builder.create(builder.getUnknownLoc()); + sourceOp->setAttr(FTD_OP_TO_SKIP, builder.getUnitAttr()); + setBBAttr(sourceOp, condBlock, builder); + + auto cstAttr = builder.getIntegerAttr(builder.getIntegerType(1), 0); + auto constOp = builder.create( + builder.getUnknownLoc(), cstAttr, sourceOp.getResult()); + constOp->setAttr(FTD_COND_VAR, builder.getUnitAttr()); + constOp->setAttr(FTD_OP_TO_SKIP, builder.getUnitAttr()); + setBBAttr(constOp, condBlock, builder); + + return constOp.getResult(); +} diff --git a/experimental/lib/Support/FtdSuppression.cpp b/experimental/lib/Support/FtdSuppression.cpp new file mode 100644 index 0000000000..b7adef506a --- /dev/null +++ b/experimental/lib/Support/FtdSuppression.cpp @@ -0,0 +1,2712 @@ +//===- FtdSuppression.cpp - Suppression Infrastructure ----------*- C++ -*-===// +// +// Dynamatic is under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// This file builds the circuitry that discards ("suppresses") a producer's +// data token on the executions where its consumer will not use it. +// +// Given the producer block and the consumer block, it reconstructs just the +// part of the control-flow graph that connects them, reduces it to the branch +// decisions that determine whether the consumer is reached, turns that into a +// Boolean condition, and emits the hardware (a tree of muxes and branches) +// that evaluates the condition and drops the token when it should not be +// delivered. +// +// Main pieces, in the order the pipeline uses them: +// - reconstruct the relevant control flow (buildLocalCFGRegion) +// - keep only the deciding branches (buildDecisionGraph) +// - analyze loops / produce an acyclic graph (CyclicGraphManager) +// - enumerate paths into a Boolean condition (enumeratePaths) +// - turn a condition into a mux-tree circuit (bddToCircuit / +// expressionToCircuit) +// - route condition tokens to the muxes (buildDistributionNetwork) +// - move condition tokens across loop levels (CyclicDemotionHelper) +// - drive the whole thing (insertDirectSuppression) +// +//===----------------------------------------------------------------------===// + +#include "experimental/Support/FtdSuppression.h" +#include "dynamatic/Analysis/ControlDependenceAnalysis.h" +#include "dynamatic/Support/Backedge.h" +#include "experimental/Support/BooleanLogic/BDD.h" +#include "experimental/Support/BooleanLogic/BoolExpression.h" +#include "experimental/Support/FtdSupport.h" +#include "mlir/Analysis/CFGLoopInfo.h" +#include "mlir/Dialect/Arith/IR/Arith.h" +#include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h" +#include "mlir/Dialect/Func/IR/FuncOps.h" +#include "mlir/Dialect/MemRef/IR/MemRef.h" +#include "mlir/IR/BuiltinOps.h" +#include "mlir/Transforms/DialectConversion.h" +using namespace mlir; +using namespace dynamatic; +using namespace dynamatic::experimental; +using namespace dynamatic::experimental::ftd; +using namespace dynamatic::experimental::boolean; + +// ===--------------------------------------------------------------------=== // +// CyclicGraphManager implementation +// ===--------------------------------------------------------------------=== // + +/// Analyzes the loop structure of `cfg`: discovers the loop-nesting scopes and +/// records each block's nesting level, so the graph can later be peeled level +/// by level. +CyclicGraphManager::CyclicGraphManager(LocalCFG &cfg) + : lcfg(cfg), domInfo(cfg.containerOp), + loopInfo(domInfo.getDomTree(cfg.region)) { + analyzeTopology(); +} + +/// Returns the loop-nesting level recorded for `bb` (0 if it is not in a loop). +unsigned CyclicGraphManager::getNestingLevel(Block *bb) const { + auto it = blockLevelMap.find(bb); + if (it != blockLevelMap.end()) + return it->second; + return 0; +} + +/// Builds the LoopScope for `loop` at nesting level `level`: records its +/// blocks, latches, and back edges, then recurses into nested loops at the next +/// level. +std::unique_ptr +CyclicGraphManager::buildScopeRecursive(mlir::CFGLoop *loop, unsigned level) { + auto scope = std::make_unique(); + scope->level = level; + scope->loopInfo = loop; + scope->header = loop->getHeader(); + + // Latches are the blocks that jump back to the header; each is a back edge. + SmallVector latches; + loop->getLoopLatches(latches); + scope->latches = latches; + + for (Block *latch : latches) { + scope->allBackEdges.insert({latch, scope->header}); + } + + auto loopBlocks = loop->getBlocks(); + for (Block *b : loopBlocks) { + scope->allBlocksInclusive.insert(b); + // Outer loops must be processed before inner loops for correct leveling. + blockLevelMap[b] = level; + } + + // Recurse into nested loops at the next level and bubble their back edges up. + for (auto *subLoop : loop->getSubLoops()) { + auto subScope = buildScopeRecursive(subLoop, level + 1); + subScope->parent = scope.get(); + for (auto &edge : subScope->allBackEdges) { + scope->allBackEdges.insert(edge); + } + scope->subLoops.push_back(std::move(subScope)); + } + return scope; +} + +/// Builds the scope tree: a level-0 top scope covering every block, plus one +/// nested scope per loop, and assigns each block its loop-nesting level. +void CyclicGraphManager::analyzeTopology() { + blockLevelMap.clear(); + // Level 0 is the acyclic top-level scope; its header stands for the entry. + topLevelScope = std::make_unique(); + topLevelScope->level = 0; + topLevelScope->header = lcfg.newProd; + topLevelScope->loopInfo = nullptr; + + // Start every block at level 0; loop members are raised while building + // scopes. + for (Block &b : lcfg.region->getBlocks()) { + topLevelScope->allBlocksInclusive.insert(&b); + blockLevelMap[&b] = 0; + } + + // One scope per top-level loop; collect their back edges into the root. + auto topLoops = loopInfo.getTopLevelLoops(); + for (auto *loop : topLoops) { + auto subScope = buildScopeRecursive(loop, 1); + subScope->parent = topLevelScope.get(); + for (auto &edge : subScope->allBackEdges) { + topLevelScope->allBackEdges.insert(edge); + } + topLevelScope->subLoops.push_back(std::move(subScope)); + } +} + +/// Produces a fresh acyclic LocalCFG for a single loop level (`scope`). The +/// scope's blocks are cloned with their back edges cut to the sink, leaving a +/// DAG: at level 0 this is the whole graph with the real consumer as the exit; +/// at a deeper level the loop header is the entry and the loop exit is the +/// consumer. Lets the rest of the pipeline reason about one loop level at a +/// time. +std::unique_ptr +CyclicGraphManager::extractLayeredCFG(const LoopScope *scope, + OpBuilder &builder) { + auto newGraph = std::make_unique(); + Location loc = builder.getUnknownLoc(); + + OpBuilder::InsertionGuard guard(builder); + auto funcType = builder.getFunctionType({}, {}); + auto dummyFunc = + builder.create(loc, "__ftd_layered_cfg__", funcType); + Region &R = dummyFunc.getBody(); + newGraph->region = &R; + newGraph->containerOp = dummyFunc; + + // Create the sink terminal (False). + // For level > 0, back-edges to own header are redirected here. + // All terminal paths eventually converge here. + Block *falseTerm = new Block(); + R.push_back(falseTerm); + newGraph->sinkBB = falseTerm; + newGraph->origMap[falseTerm] = nullptr; + + // Create the consumer terminal (True). + // For level > 0 this represents the loop exit. + // For level 0 this maps to the actual consumer in the original CFG. + Block *trueTerm = new Block(); + R.push_back(trueTerm); + newGraph->newCons = trueTerm; + + if (scope->level == 0) + newGraph->origMap[trueTerm] = lcfg.origMap.lookup(lcfg.newCons); + else + newGraph->origMap[trueTerm] = nullptr; + + // Clone the header block as the entry (Producer) of this layer. + Block *clonedHeader = new Block(); + R.push_back(clonedHeader); + newGraph->newProd = clonedHeader; + newGraph->origMap[clonedHeader] = lcfg.origMap.lookup(scope->header); + + // Build a mapping from original blocks to their cloned counterparts. + DenseMap clonedMap; + clonedMap[scope->header] = clonedHeader; + + // Build topo position map from the input graph for residual back-edge + // detection (handles irreducible cycles not captured by CFGLoopInfo). + DenseMap inputTopoPos; + for (unsigned i = 0; i < lcfg.topoOrder.size(); ++i) + inputTopoPos[lcfg.topoOrder[i]] = i; + + // At level 0, map the original consumer and sink to the terminals + // so that in-scope edges targeting them are resolved correctly. + if (scope->level == 0) { + clonedMap[lcfg.newCons] = trueTerm; + clonedMap[lcfg.sinkBB] = falseTerm; + } + + // Clone all in-scope blocks except those already handled above. + for (Block *b : scope->allBlocksInclusive) { + if (b == scope->header || b == lcfg.newCons || b == lcfg.sinkBB) + continue; + // Skip the artificial dummyStart block created by buildDecisionGraph. + // It has no predecessors (it IS the region entry), unconditionally + // branches to scope->header (newProd), and carries no decision logic. + // Cloning it would create an orphan that ends up as region.front(), + // corrupting downstream CDA which uses region.front() as its entry. + if (b->hasNoPredecessors() && b != scope->header) + continue; + Block *nb = new Block(); + R.push_back(nb); + clonedMap[b] = nb; + newGraph->origMap[nb] = lcfg.origMap.lookup(b); + } + + // Reconstruct edges and terminators for each cloned block. + for (auto [origBlock, newBlock] : clonedMap) { + // Skip terminal blocks; they receive their terminators below. + if (origBlock == lcfg.newCons || origBlock == lcfg.sinkBB) + continue; + + builder.setInsertionPointToEnd(newBlock); + Operation *origTerm = origBlock->getTerminator(); + + if (!origTerm) { + llvm::errs() << "Warning: Block without terminator found in LocalCFG " + "extraction.\n"; + continue; + } + + SmallVector validSuccessors; + SmallVector keepSuccessor; + + for (unsigned i = 0; i < origTerm->getNumSuccessors(); ++i) { + Block *origSucc = origTerm->getSuccessor(i); + + if (origSucc == scope->header) { + // Back-edge to this scope's own header. + // Level 0: header is dummystart; no real back-edge should exist, + // but cut defensively if encountered. + // Level > 0: redirect to the sink terminal. + if (scope->level == 0) { + validSuccessors.push_back(nullptr); + keepSuccessor.push_back(false); + } else { + validSuccessors.push_back(falseTerm); + keepSuccessor.push_back(true); + } + } else if (scope->allBackEdges.contains({origBlock, origSucc})) { + // Deep back-edge belonging to an inner loop. + // Always cut; the inner loop's own layer handles it. + validSuccessors.push_back(nullptr); + keepSuccessor.push_back(false); + } else if (clonedMap.count(origSucc)) { + // Standard in-scope edge. + // Check for residual back-edges (irreducible cycles) using the + // input graph's topological order. These are not natural loop + // back-edges and thus not in allBackEdges, but still form cycles. + if (inputTopoPos.count(origSucc) && inputTopoPos.count(origBlock) && + inputTopoPos[origSucc] <= inputTopoPos[origBlock]) { + // Residual back-edge — redirect to sink. + validSuccessors.push_back(falseTerm); + keepSuccessor.push_back(true); + } else { + validSuccessors.push_back(clonedMap[origSucc]); + keepSuccessor.push_back(true); + } + } else { + // Edge leaving the current scope. + if (scope->level == 0) { + llvm::errs() + << "[FTD Warning] Block outside Level 0 scope encountered.\n"; + } + // Redirect to the consumer terminal (represents loop exit). + validSuccessors.push_back(trueTerm); + keepSuccessor.push_back(true); + } + } + + // Rebuild the terminator based on surviving successors. + if (auto cbr = dyn_cast(origTerm)) { + if (keepSuccessor[0] && keepSuccessor[1]) { + // Both paths survived; keep conditional branch. + Value cond = builder.create(loc, 1, 1); + builder.create(loc, cond, validSuccessors[0], + validSuccessors[1]); + } else if (keepSuccessor[0] && !keepSuccessor[1]) { + // Only the true path survived. + builder.create(loc, validSuccessors[0]); + } else if (!keepSuccessor[0] && keepSuccessor[1]) { + // Only the false path survived. + builder.create(loc, validSuccessors[1]); + } + // If neither survived the block is left without a terminator (dead). + } else if (auto br = dyn_cast(origTerm)) { + if (keepSuccessor[0]) { + builder.create(loc, validSuccessors[0]); + } + // If the sole successor was cut, no terminator is created (dead). + } + } + + builder.setInsertionPointToEnd(trueTerm); + builder.create(loc, falseTerm); + builder.setInsertionPointToEnd(falseTerm); + builder.create(loc); + + // Graph simplification. + // + // Level 0: only remove dead blocks (those without a terminator that are + // not terminals). This eliminates unreachable subgraphs created by + // cutting deep back-edges. No duplicate-successor merging and no path + // compression are performed; those are deferred to the caller. + // + // Level > 0: full canonicalization. After this pass every non-terminal + // block should carry a conditional branch, which is what the BDD builder + // expects. + bool changed = true; + while (changed) { + changed = false; + + for (Block &block : llvm::make_early_inc_range(R)) { + // Never touch the entry or the two terminal blocks. + if (&block == newGraph->newProd || &block == newGraph->newCons || + &block == newGraph->sinkBB) + continue; + + Operation *term = block.getTerminator(); + + // Dead block removal (applies to all levels). + // A block without a terminator has no outgoing edges and is dead. + // Disconnect its predecessors so the removal can propagate. + if (!term) { + SmallVector predTerms; + DenseSet predTermSeen; + for (auto &use : block.getUses()) + if (predTermSeen.insert(use.getOwner()).second) + predTerms.push_back(use.getOwner()); + + for (Operation *predTerm : predTerms) { + OpBuilder localBuilder(predTerm->getContext()); + localBuilder.setInsertionPoint(predTerm); + + if (auto br = dyn_cast(predTerm)) { + // Predecessor unconditionally jumps here; it becomes dead too. + predTerm->erase(); + } else if (auto cbr = dyn_cast(predTerm)) { + Block *trueDest = cbr.getTrueDest(); + Block *falseDest = cbr.getFalseDest(); + + if (trueDest == &block && falseDest == &block) { + // Both legs land on this dead block; predecessor dies. + predTerm->erase(); + } else if (trueDest == &block) { + // True leg is dead; degrade to unconditional false branch. + localBuilder.create(loc, falseDest); + predTerm->erase(); + } else if (falseDest == &block) { + // False leg is dead; degrade to unconditional true branch. + localBuilder.create(loc, trueDest); + predTerm->erase(); + } + } + } + + block.erase(); + changed = true; + continue; + } + + // The remaining optimizations only apply to level > 0. + if (scope->level == 0) + continue; + + // Merge duplicate successors: CondBranch(A, A) becomes Branch(A). + if (auto condBr = dyn_cast(term)) { + if (condBr.getTrueDest() == condBr.getFalseDest()) { + OpBuilder localBuilder(term->getContext()); + localBuilder.setInsertionPoint(term); + localBuilder.create(loc, condBr.getTrueDest()); + term->erase(); + term = block.getTerminator(); + changed = true; + } + } + + // Path compression: if a block unconditionally jumps to a single + // destination, redirect all predecessors directly to that destination + // and remove the block. + if (auto br = dyn_cast(term)) { + Block *dest = br.getDest(); + block.replaceAllUsesWith(dest); + block.erase(); + changed = true; + continue; + } + } + } + + // Compute topological order starting from the producer. + DenseSet visited; + std::function topo = [&](Block *u) { + if (!u || visited.contains(u)) + return; + visited.insert(u); + if (auto *term = u->getTerminator()) + for (auto it = term->successor_begin(); it != term->successor_end(); ++it) + topo(*it); + newGraph->topoOrder.push_back(u); + }; + + topo(newGraph->newProd); + std::reverse(newGraph->topoOrder.begin(), newGraph->topoOrder.end()); + + // Erase orphan blocks: any block not reachable from newProd (i.e. not in + // the visited set) is structural debris. If left in the region it may + // end up before newProd in the block list, becoming region.front() and + // corrupting downstream analyses (CDA, dominance) that assume + // region.front() is the entry. + for (Block &block : llvm::make_early_inc_range(R)) { + if (!visited.contains(&block)) { + // Erase the terminator first to drop outgoing BlockOperand refs. + if (Operation *term = block.getTerminator()) + term->erase(); + // The orphan is unreachable, so no other terminator references it. + block.erase(); + } + } + + // Physical reordering: ensure region.front() == newProd so that + // downstream CDA (which starts DFS from region.front()) uses the + // correct entry block. + for (Block *b : newGraph->topoOrder) { + if (b != newGraph->sinkBB) { + b->moveBefore(newGraph->sinkBB); + } + } + + return newGraph; +} + +// ===--------------------------------------------------------------------=== // +// SignalRegistry implementation +// ===--------------------------------------------------------------------=== // + +/// Records `val` as the signal for `var` produced on control-flow path `path`. +void SignalRegistry::registerSignal(StringRef var, const PathContext &path, + Value val) { + map[var.str()].push_back({path, val}); +} + +/// Returns the signal for `var` whose registered path is the longest prefix of +/// `queryPath` (the closest definition on the current path), or null if none. +Value SignalRegistry::lookup(StringRef var, const PathContext &queryPath) { + std::string v = var.str(); + if (map.find(v) == map.end()) + return nullptr; + + Value bestMatch = nullptr; + size_t bestLen = 0; + bool foundAny = false; + + for (auto &entry : map[v]) { + const PathContext ®Path = entry.first; + if (regPath.size() > queryPath.size()) + continue; + + // Filter: The registered path must be a prefix of the query path + // to ensure the signal lies on the same control flow path. + bool isPrefix = true; + for (size_t i = 0; i < regPath.size(); ++i) { + if (regPath[i] != queryPath[i]) { + isPrefix = false; + break; + } + } + + // Selection: Choose the longest matching prefix (closest definition). + if (isPrefix) { + if (!foundAny || regPath.size() >= bestLen) { + bestLen = regPath.size(); + bestMatch = entry.second; + foundAny = true; + } + } + } + return bestMatch; +} + +// ===--------------------------------------------------------------------=== // +// Static helpers (internal to this TU) +// ===--------------------------------------------------------------------=== // + +/// Identify the block that has muxCondition as its terminator condition +/// Note that it is not necessarily the same block defining the muxCondition +static Block *returnMuxConditionBlock(Value muxCondition, + ftd::ShadowCFG &shadow) { + while (true) { + Operation *defOp = muxCondition.getDefiningOp(); + if (!defOp) + return shadow.getBlock(0); + + // Found the NotIOp placeholder — read its BB attribute. + if (defOp->hasAttr(FTD_COND_VAR)) { + auto bbAttr = defOp->getAttrOfType("handshake.bb"); + assert(bbAttr && "Condition placeholder missing handshake.bb"); + return shadow.getBlock(bbAttr.getUInt()); + } + + // Trace backward through suppression branches (marked FTD_OP_TO_SKIP). + if (isa(defOp) && + defOp->hasAttr(FTD_OP_TO_SKIP)) { + muxCondition = defOp->getOperand(1); + continue; + } + + // Fallback: use handshake.bb on the defining op. + if (auto bbAttr = defOp->getAttrOfType("handshake.bb")) + return shadow.getBlock(bbAttr.getUInt()); + + return shadow.getBlock(0); + } +} + +/// When the consumer is a loop header that contains the producer, suppression +/// circuitry is better placed at the loop's exit than at the producer's block. +/// Returns the handshake.bb attribute of that first exit block, or an empty +/// attribute when this situation does not apply. +static IntegerAttr getFirstLoopExitBBAttrIfHeaderConsumer( + OpBuilder &builder, Region ®ion, Block *producerBlock, + Block *consumerBlock, const ftd::BlockIndexing &bi, + ftd::ShadowCFG &shadow) { + if (!producerBlock || !consumerBlock || producerBlock == consumerBlock) + return {}; + + DominanceInfo domInfo; + mlir::CFGLoopInfo loopInfo(domInfo.getDomTree(®ion)); + CFGLoop *consumerLoop = loopInfo.getLoopFor(consumerBlock); + if (!consumerLoop || consumerLoop->getHeader() != consumerBlock || + !consumerLoop->contains(producerBlock)) + return {}; + + SmallVector exitBlocks; + consumerLoop->getExitBlocks(exitBlocks); + if (exitBlocks.empty()) + return {}; + + llvm::sort(exitBlocks, + [&](Block *lhs, Block *rhs) { return bi.isLess(lhs, rhs); }); + + unsigned exitBBIdx = shadow.getBlockIndex(exitBlocks.front()); + return ftd::getBBIndexAttr(builder.getContext(), exitBBIdx); +} + +/// Helper function to generate expression combining Local Logic (True/False) +/// with Original Variables (c2, c3...). +/// Note: Input is Local Deps +static boolean::BoolExpression * +getHybridPathExpression(const std::vector &localPath, + const ftd::LocalCFG &lcfg, const ftd::BlockIndexing &bi, + const DenseSet &localDeps) { + + // Start with 1 + auto *exp = boolean::BoolExpression::boolOne(); + + unsigned pathSize = localPath.size(); + for (unsigned i = 0; i < pathSize - 1; i++) { + // Local Block + Block *u = localPath[i]; + // Local Block (Target) + Block *v = localPath[i + 1]; + + // 1. Check Dependency using LOCAL sets + if (localDeps.contains(u)) { + Operation *term = u->getTerminator(); + + if (isa(term)) { + // 2. Map to Original Block to get the Variable Index/Name + Block *origU = lcfg.origMap.lookup(u); + + // Special handling for second visit if needed, though usually origMap + // covers it + if (!origU && u == lcfg.secondVisitBB) { + origU = lcfg.origMap.lookup(lcfg.newProd); + } + + if (origU) { + // 3. Get Name from ORIGINAL CFG (e.g., "c2") + auto blockIndexOptional = bi.getIndexFromBlock(origU); + if (blockIndexOptional.has_value()) { + std::string blockCondition = bi.getBlockCondition(origU); + boolean::BoolExpression *cond = + boolean::BoolExpression::parseSop(blockCondition); + + // 4. Get Polarity from LOCAL CFG (Structure) + // Does the path go to 'v' via the False branch in the Decision + // Graph? + auto condOp = dyn_cast(term); + if (condOp.getFalseDest() == v) { + // Add '~' if Local Graph says False + cond->boolNegate(); + } + + // Combine + exp = boolean::BoolExpression::boolAnd(exp, cond); + } + } + } + } + } + return exp; +} + +// ===--------------------------------------------------------------------=== // +// Path enumeration on LocalCFG +// ===--------------------------------------------------------------------=== // + +/// Enumerates every producer-to-consumer path in `lcfg` and ORs their path +/// conditions into one expression: the condition under which the consumer is +/// reached. +BoolExpression *ftd::enumeratePaths(const ftd::LocalCFG &lcfg, + const ftd::BlockIndexing &bi, + const DenseSet &controlDeps) { + + // 1. Path Finding using Iterative DFS (on Local CFG) + std::vector> allPaths; + + struct StackFrame { + Block *u; + unsigned currIdx; + unsigned numSuccs; + }; + + std::vector dfsStack; + std::vector currentLocalPath; + + if (lcfg.newProd && lcfg.newCons) { + Block *root = lcfg.newProd; + auto *term = root->getTerminator(); + unsigned n = term ? term->getNumSuccessors() : 0; + + dfsStack.push_back({root, 0, n}); + currentLocalPath.push_back(root); + } else { + return BoolExpression::boolZero(); + } + + while (!dfsStack.empty()) { + StackFrame &frame = dfsStack.back(); + + // Case A: Reached Consumer + if (frame.u == lcfg.newCons) { + // [CRITICAL] Store the LOCAL path exactly as traversed. + // Do NOT map to original blocks here. + allPaths.push_back(currentLocalPath); + + currentLocalPath.pop_back(); + dfsStack.pop_back(); + continue; + } + + // Case B: Traverse Successors + if (frame.currIdx < frame.numSuccs) { + auto *term = frame.u->getTerminator(); + Block *succ = term->getSuccessor(frame.currIdx); + frame.currIdx++; + + bool isCycle = std::find(currentLocalPath.begin(), currentLocalPath.end(), + succ) != currentLocalPath.end(); + + if (succ != lcfg.sinkBB && !isCycle) { + auto *succTerm = succ->getTerminator(); + unsigned succN = succTerm ? succTerm->getNumSuccessors() : 0; + + dfsStack.push_back({succ, 0, succN}); + currentLocalPath.push_back(succ); + } + } + // Case C: Backtrack + else { + currentLocalPath.pop_back(); + dfsStack.pop_back(); + } + } + + if (allPaths.empty()) + return BoolExpression::boolZero(); + + // 2. Expression Generation + BoolExpression *sop = BoolExpression::boolZero(); + + for (const std::vector &path : allPaths) { + // Use the hybrid helper to look up logic locally and names globally + BoolExpression *minterm = + getHybridPathExpression(path, lcfg, bi, controlDeps); + + sop = BoolExpression::boolOr(sop, minterm); + } + return sop->boolMinimizeSop(); +} + +// ===--------------------------------------------------------------------=== // +// BDD → Circuit conversion +// ===--------------------------------------------------------------------=== // + +/// Retrieves the initial value from BlockIndexing. +static Value +getOriginalValue(mlir::OpBuilder &builder, StringRef varName, + const ftd::BlockIndexing &bi, + DenseMap> *pendingMuxOperands, + ftd::ShadowCFG *shadow = nullptr) { + StringRef lookupName = varName; + if (lookupName.startswith("~")) { + llvm::errs() << "[FTD Error] Negated variable '" << varName << "'.\n"; + lookupName = lookupName.drop_front(); + } + + auto conditionOpt = bi.getBlockFromCondition(lookupName.str()); + if (!conditionOpt.has_value()) + return nullptr; + + Value condition; + if (shadow) { + // conditionOpt is the shadow Block* looked up by BlockIndexing. + // Convert it to enumeration index for conditionMap lookup. + Block *shadowBlock = conditionOpt.value(); + unsigned enumIdx = shadow->getBlockIndex(shadowBlock); + condition = shadow->getCondition(enumIdx); + } else { + // Pre-flatten CfToHandshake path: use a SourceOp placeholder instead of + // the real terminator condition. This avoids creating backedges for + // condition values (which caused null-Value crashes in bddToCircuit). + condition = getOrCreateCondPlaceholder(conditionOpt.value(), builder); + } + + return condition; +} + +/// Converts a boolean expression node to a circuit signal. +static Value boolExpressionToCircuit( + mlir::OpBuilder &builder, experimental::boolean::BoolExpression *expr, + Block *block, SignalRegistry ®istry, const PathContext ¤tPath, + const ftd::BlockIndexing &bi, + DenseMap> *pendingMuxOperands, + ftd::ShadowCFG *shadow = nullptr, IntegerAttr forcedBBAttr = {}) { + + // Case 1: Variable + if (expr->type == ExpressionType::Variable) { + SingleCond *singleCond = static_cast(expr); + std::string varName = singleCond->id; + + // 1. Registry Lookup + Value val = registry.lookup(varName, currentPath); + + // 2. Fallback + if (!val) { + val = getOriginalValue(builder, varName, bi, pendingMuxOperands, shadow); + + if (!val) { + llvm::errs() << "[FTD Error] Variable '" << varName + << "' not found in Registry or BlockIndexing.\n"; + assert(val && "Signal missing from IR"); + } + + if (!val.getType().isa()) { + val.setType(ftd::channelifyType(val.getType())); + } + } + + // 3. Handle Negation + if (singleCond->isNegated) { + builder.setInsertionPointToStart(block); + Location loc = block->getOperations().front().getLoc(); + Type chanI1 = ftd::channelifyType(builder.getI1Type()); + + auto notOp = builder.create(loc, chanI1, val); + notOp->setAttr(FTD_OP_TO_SKIP, builder.getUnitAttr()); + setBBAttrWithFallback(notOp, forcedBBAttr, block, builder); + return notOp->getResult(0); + } + + return val; + } + + // Case 2: Constant + auto sourceOp = builder.create(block->front().getLoc()); + auto intType = builder.getIntegerType(1); + int constVal = (expr->type == ExpressionType::One ? 1 : 0); + auto cstAttr = builder.getIntegerAttr(intType, constVal); + auto constOp = builder.create( + block->front().getLoc(), cstAttr, sourceOp.getResult()); + constOp->setAttr(FTD_OP_TO_SKIP, builder.getUnitAttr()); + setBBAttrWithFallback(sourceOp, forcedBBAttr, block, builder); + setBBAttrWithFallback(constOp, forcedBBAttr, block, builder); + + return constOp.getResult(); +} + +/// Recursively converts a BDD to a Mux Tree. +Value ftd::bddToCircuit( + mlir::OpBuilder &builder, BDD *bdd, Block *block, SignalRegistry ®istry, + PathContext currentPath, const ftd::BlockIndexing &bi, + DenseMap> *pendingMuxOperands, + ftd::ShadowCFG *shadow, IntegerAttr forcedBBAttr) { + using namespace experimental::boolean; + + // 1. Leaf Node + if (!bdd->successors.has_value()) { + return boolExpressionToCircuit(builder, bdd->boolVariable, block, registry, + currentPath, bi, pendingMuxOperands, shadow, + forcedBBAttr); + } + + // 2. Mux Node + std::string varName = bdd->boolVariable->toString(); + + Value muxCond = registry.lookup(varName, currentPath); + if (!muxCond) { + muxCond = + getOriginalValue(builder, varName, bi, pendingMuxOperands, shadow); + assert(muxCond && "Mux condition not found"); + if (!muxCond.getType().isa()) + muxCond.setType(ftd::channelifyType(muxCond.getType())); + } + + SmallVector muxOperands; + + // Recursion: Update PathContext so downstream lookups find distributed + // signals + PathContext falsePath = currentPath; + falsePath.push_back({varName, false}); + muxOperands.push_back(bddToCircuit(builder, bdd->successors.value().first, + block, registry, falsePath, bi, + pendingMuxOperands, shadow, forcedBBAttr)); + + PathContext truePath = currentPath; + truePath.push_back({varName, true}); + muxOperands.push_back(bddToCircuit(builder, bdd->successors.value().second, + block, registry, truePath, bi, + pendingMuxOperands, shadow, forcedBBAttr)); + + auto muxOp = builder.create( + muxCond.getLoc(), muxOperands[0].getType(), muxCond, muxOperands); + muxOp->setAttr(FTD_OP_TO_SKIP, builder.getUnitAttr()); + setBBAttrWithFallback(muxOp, forcedBBAttr, block, builder); + + return muxOp.getResult(); +} + +// ===--------------------------------------------------------------------=== // +// Token distribution +// ===--------------------------------------------------------------------=== // + +/// Generates the Suppression Logic (Mux Tree) for a branch's select signal. +/// It constructs the logic for the "UNREACHABLE" condition. +/// F_suppress = NOT( OR( All Valid Paths ) ) +static Value generateReachabilityLogic( + mlir::OpBuilder &builder, Block *block, + const std::vector &requirements, + const PathContext ¤tPath, SignalRegistry ®istry, + const ftd::BlockIndexing &bi, size_t startIndex, + DenseMap> *pendingMuxOperands, + ftd::ShadowCFG *shadow = nullptr) { + + using namespace experimental::boolean; + + // 1. Construct Boolean Expression for Valid Paths + BoolExpression *fValid = BoolExpression::boolZero(); + + for (const auto &req : requirements) { + BoolExpression *pathExpr = BoolExpression::boolOne(); + + // Iterate through the path suffix starting from the current split point + for (size_t i = startIndex; i < req.path.size(); ++i) { + PathStep step = req.path[i]; + // Construct: SingleCond(Type, Name, Negated) + BoolExpression *stepExpr = + new SingleCond(ExpressionType::Variable, step.var, !step.value); + pathExpr = BoolExpression::boolAnd(pathExpr, stepExpr); + } + fValid = BoolExpression::boolOr(fValid, pathExpr); + } + + // 2. Compute Suppression Condition: F_suppress = NOT( F_valid ) + // We want the circuit to output TRUE when the path is INVALID. + BoolExpression *fSuppress = fValid->boolNegate(); + fSuppress = fSuppress->boolMinimize(); + + // 3. Build BDD and Circuit for Suppression Condition + std::set vars = fSuppress->getVariables(); + std::vector cofactorList(vars.begin(), vars.end()); + + // Sort cofactor list to match topological order (e.g., c0, then c1) + // This ensures the Mux Tree structure matches the dependency order. + std::sort(cofactorList.begin(), cofactorList.end(), + [&](const std::string &a, const std::string &b) { + auto idA = bi.getBlockFromCondition(a); + auto idB = bi.getBlockFromCondition(b); + if (!idA || !idB) + return a < b; + return bi.isLess(idA.value(), idB.value()); + }); + + BDD *bdd = buildBDD(fSuppress, cofactorList); + + // Note: bddToCircuit uses registry.lookup. If the suppression logic involves + // variables distributed earlier (like c3a), it will correctly find them. + return bddToCircuit(builder, bdd, block, registry, currentPath, bi, + pendingMuxOperands, shadow); +} + +/// Recursively builds the Branch Tree. +static void buildBranchTreeRecursive( + mlir::OpBuilder &builder, StringRef currentVar, + std::vector &requirements, PathContext currentPath, + SignalRegistry ®istry, const ftd::BlockIndexing &bi, + DenseMap> *pendingMuxOperands, + ftd::ShadowCFG *shadow = nullptr) { + + // 1. Retrieve Data Signal + // Look up the current data signal to be distributed from the registry using + // the current path context. + Value sourceVal = registry.lookup(currentVar, currentPath); + assert(sourceVal && "Source value for distribution not found"); + + // 2. Identify Split Variable + // Key: {Variable Name, Value} -> Value: List of requirements. + std::map, std::vector> + groups; + std::string splitVar = ""; + bool splitFound = false; + + size_t maxDepth = 0; + for (const auto &req : requirements) + maxDepth = std::max(maxDepth, req.path.size()); + + // Scan forward to find the first point where requirements disagree on a + // variable value. + size_t scanDepth = currentPath.size(); + for (; scanDepth < maxDepth; ++scanDepth) { + groups.clear(); + splitVar = ""; + + // Simply collect the step at this depth. + for (auto &req : requirements) { + PathStep step = req.path[scanDepth]; + if (splitVar == "") + splitVar = step.var; + + groups[{step.var, step.value}].push_back(req); + } + + // Divergence found. + if (groups.size() > 1) { + splitFound = true; + break; + } + } + + if (!splitFound) + return; + + // 3. Retrieve Raw Select Signal + // We need the physical control signal corresponding to the 'splitVar' found + // above. + Value conditionVal = registry.lookup(splitVar, currentPath); + if (!conditionVal) { + // Fallback: If not in registry, get the original value from the IR + // (BlockIndexing). + conditionVal = + getOriginalValue(builder, splitVar, bi, pendingMuxOperands, shadow); + } + assert(conditionVal && "Splitter condition value not found"); + + // Ensure Types are compatible with Handshake channels. + if (!conditionVal.getType().isa()) + conditionVal.setType(ftd::channelifyType(conditionVal.getType())); + if (!sourceVal.getType().isa()) + sourceVal.setType(ftd::channelifyType(sourceVal.getType())); + + // 4. Register Outputs and Recurse + // [Context Backfilling] + // Since we might have skipped several variables (Common Prefix) to reach + // 'splitVar', we must fill these skipped steps back into the PathContext. + // This ensures that the recursive call has a continuous path history, keeping + // index alignment correct. + PathContext baseNextPath = currentPath; + if (!groups.empty()) { + // Take the first requirement as a template to retrieve the skipped steps. + const auto &repReq = groups.begin()->second.front(); + for (size_t k = currentPath.size(); k < scanDepth; ++k) + baseNextPath.push_back({repReq.path[k].var, repReq.path[k].value}); + } + + // 5. [Suppression Logic] + // Generate the logic to identify "Unreachable" or "Invalid" paths. + // We pass 'scanDepth' (the index of splitVar) to the generator. + // This tells the generator to check validity starting from the current split + // variable, effectively ignoring the "Common Prefix" variables skipped in + // Step 2 (which are implicitly valid). The logic checks the entire future + // path to ensure reachability. + Value suppressCondition = generateReachabilityLogic( + builder, sourceVal.getParentBlock(), requirements, baseNextPath, registry, + bi, scanDepth, pendingMuxOperands, shadow); + + if (!suppressCondition.getType().isa()) + suppressCondition.setType(ftd::channelifyType(suppressCondition.getType())); + + // [Suppression Branch] + // Acts as a filter: + // If suppressCondition is TRUE (Invalid Path) -> Output to Sink (Discard + // Token). If suppressCondition is FALSE (Valid Path) -> Output to + // 'activeSelectSignal' (Pass Token). + SmallVector suppResultTypes = {conditionVal.getType(), + conditionVal.getType()}; + auto suppBranch = builder.create( + conditionVal.getLoc(), suppResultTypes, suppressCondition, conditionVal); + suppBranch->setAttr(FTD_OP_TO_SKIP, builder.getUnitAttr()); + setBBAttr(suppBranch, conditionVal.getParentBlock(), builder); + + // False Output -> Active Select (Pass to Main Branch) + Value activeSelectSignal = suppBranch.getFalseResult(); + + // 6. [Distribution Logic] Main Branch + SmallVector resultTypes = {sourceVal.getType(), sourceVal.getType()}; + + // Create the branch that splits the 'sourceVal' based on the (possibly + // filtered) 'activeSelectSignal'. + auto branchOp = builder.create( + sourceVal.getLoc(), resultTypes, activeSelectSignal, sourceVal); + branchOp->setAttr(FTD_OP_TO_SKIP, builder.getUnitAttr()); + setBBAttr(branchOp, sourceVal.getParentBlock(), builder); + + Value trueResult = branchOp.getTrueResult(); + Value falseResult = branchOp.getFalseResult(); + + // Handle the True branch recursion + if (!groups[{splitVar, true}].empty()) { + PathContext truePath = baseNextPath; + truePath.push_back({splitVar, true}); + registry.registerSignal(currentVar, truePath, trueResult); + buildBranchTreeRecursive(builder, currentVar, groups[{splitVar, true}], + truePath, registry, bi, pendingMuxOperands, + shadow); + } + + // Handle the False branch recursion + if (!groups[{splitVar, false}].empty()) { + PathContext falsePath = baseNextPath; + falsePath.push_back({splitVar, false}); + registry.registerSignal(currentVar, falsePath, falseResult); + buildBranchTreeRecursive(builder, currentVar, groups[{splitVar, false}], + falsePath, registry, bi, pendingMuxOperands, + shadow); + } +} + +/// Main entry point of distribution logic. +void ftd::buildDistributionNetwork( + mlir::OpBuilder &builder, const ftd::LocalCFG &lcfg, + const ftd::BlockIndexing &bi, SignalRegistry ®istry, + DenseMap> *pendingMuxOperands, + ftd::ShadowCFG *shadow) { + using namespace experimental::boolean; + + // 1. Collect Variable Requirements + std::map> varNeeds; + DenseSet collectOnStack; + std::function collect = [&](Block *curr, + PathContext path) { + // Stop recursion if reaching the consumer or the sink block + if (curr == lcfg.newCons || curr == lcfg.sinkBB) + return; + + // Cycle detection: stop if this block is already on the current DFS path + if (!collectOnStack.insert(curr).second) + return; + + // Find the condition variable + Block *origBlock = lcfg.origMap.lookup(curr); + std::string var = ""; + if (origBlock) + var = bi.getBlockCondition(origBlock); + + // Record the variable requirement + if (!var.empty()) { + varNeeds[var].push_back({var, path}); + } + + auto *term = curr->getTerminator(); + if (!term) + return; + + // Handle conditional branches. + if (auto condBr = dyn_cast(term)) { + if (!var.empty()) { + PathContext truePath = path; + truePath.push_back({var, true}); + collect(condBr.getTrueDest(), truePath); + + PathContext falsePath = path; + falsePath.push_back({var, false}); + collect(condBr.getFalseDest(), falsePath); + } else { + llvm::errs() << "[FTD ERROR] CondBranchOp encountered with empty " + "condition variable at block " + << origBlock << ". Successors will not be traversed.\n"; + } + } else { + // Handle unconditional branches + for (Block *succ : term->getSuccessors()) { + collect(succ, path); + } + } + + collectOnStack.erase(curr); + }; + + // Start collection from the producer block + if (lcfg.newProd) { + collect(lcfg.newProd, {}); + } + + // 2. Topological Sort + std::vector sortedVars; + for (auto &kv : varNeeds) + sortedVars.push_back(kv.first); + + std::sort(sortedVars.begin(), sortedVars.end(), + [&](const std::string &a, const std::string &b) { + auto idA = bi.getBlockFromCondition(a); + auto idB = bi.getBlockFromCondition(b); + if (!idA || !idB) { + llvm::errs() + << "[FTD Warning] Variable missing from BlockIndexing: '" + << (idA ? b : a) << "'\n"; + return a < b; + } + return bi.isLess(idA.value(), idB.value()); + }); + + // 3. Initial Registration and Construct Branch Trees + for (const auto &var : sortedVars) { + // Use pre-registered value if available (e.g. demoted high-level variable) + Value rawVal = registry.lookup(var, {}); + if (!rawVal) { + rawVal = getOriginalValue(builder, var, bi, pendingMuxOperands, shadow); + if (!rawVal) { + llvm::errs() << "[FTD Error] Variable '" << var + << "' not found in BlockIndexing during registration.\n"; + assert(rawVal && "Signal missing from IR"); + } + if (!rawVal.getType().isa()) + rawVal.setType(ftd::channelifyType(rawVal.getType())); + registry.registerSignal(var, {}, rawVal); + } + if (varNeeds[var].size() > 1) { + buildBranchTreeRecursive(builder, var, varNeeds[var], {}, registry, bi, + pendingMuxOperands, shadow); + } + } +} + +// ===--------------------------------------------------------------------=== // +// expressionToCircuit — unified BDD pipeline +// ===--------------------------------------------------------------------=== // + +/// Numbers each original block by its position in `lcfg`'s topological order, +/// giving the variable order used when lowering an expression to a mux tree. +DenseMap ftd::computeTopoRank(const LocalCFG &lcfg) { + DenseMap rank; + unsigned i = 0; + // Number each original block by its position in topological order. + for (Block *b : lcfg.topoOrder) + if (auto *ob = lcfg.origMap.lookup(b)) + rank[ob] = i++; + return rank; +} + +/// Lowers a Boolean expression to a mux-tree circuit: minimize it, order its +/// variables by `varRank`, build the BDD, and emit the muxes. +Value ftd::expressionToCircuit( + OpBuilder &builder, BoolExpression *expr, + const DenseMap &varRank, Block *insertBlock, + SignalRegistry ®istry, const BlockIndexing &bi, + DenseMap> *pendingMuxOperands, + ShadowCFG *shadow, IntegerAttr forcedBBAttr) { + + // Minimize, order the variables by block position (the BDD variable order), + // build the BDD, and lower it to a mux tree. + expr = expr->boolMinimize(); + std::set vars = expr->getVariables(); + + // Sort cofactors by topological rank + std::vector> tmp; + tmp.reserve(vars.size()); + for (auto &var : vars) + if (auto blkOpt = bi.getBlockFromCondition(var)) + if (varRank.count(*blkOpt)) + tmp.emplace_back(varRank.lookup(*blkOpt), var); + llvm::sort(tmp, [](auto &a, auto &b) { return a.first < b.first; }); + + std::vector cofactorList; + cofactorList.reserve(tmp.size()); + for (auto &p : tmp) + cofactorList.push_back(p.second); + + BDD *bdd = buildBDD(expr, cofactorList); + return bddToCircuit(builder, bdd, insertBlock, registry, {}, bi, + pendingMuxOperands, shadow, forcedBBAttr); +} + +/// Same as above, deriving the variable order from `rankSource`'s topology. +Value ftd::expressionToCircuit( + OpBuilder &builder, BoolExpression *expr, const LocalCFG &rankSource, + Block *insertBlock, SignalRegistry ®istry, const BlockIndexing &bi, + DenseMap> *pendingMuxOperands, + ShadowCFG *shadow, IntegerAttr forcedBBAttr) { + DenseMap rank = computeTopoRank(rankSource); + return expressionToCircuit(builder, expr, rank, insertBlock, registry, bi, + pendingMuxOperands, shadow, forcedBBAttr); +} + +// ===--------------------------------------------------------------------=== // +// computeLoopBackedgeCondition — shared MU/regen pipeline +// ===--------------------------------------------------------------------=== // + +/// Computes the condition under which the loop at `loopHeader` re-enters the +/// header (takes its back edge). Builds a self-loop local CFG (producer = +/// consumer = header), reduces it to level 0, routes the condition tokens, and +/// lowers the resulting expression to a circuit. Used by the MU/regen path. +Value ftd::computeLoopBackedgeCondition( + OpBuilder &builder, Block *loopHeader, Block *insertBlock, + const BlockIndexing &bi, + DenseMap> *pendingMuxOperands, + ShadowCFG *shadow) { + + // 1. Build Local CFG with Prod = Cons = Loop Header (self-loop graph) + OpBuilder tmpBuilder(builder.getContext()); + auto locGraph = buildLocalCFGRegion(tmpBuilder, loopHeader, loopHeader, bi); + + // 2. CDA on raw graph to get full control deps + ControlDependenceAnalysis locCDATmp(*locGraph->region); + DenseSet locConsControlDepsFull = + locCDATmp.getAllBlockDeps()[locGraph->newCons].allControlDeps; + DenseSet locConsAcyclicDeps = + locCDATmp.getAllBlockDeps()[locGraph->newCons].allControlDeps; + + // 3. Build full decision graph for CyclicGraphManager + auto fullDecisionGraph = + buildDecisionGraph(*locGraph, locConsControlDepsFull); + CyclicGraphManager cyclicMgr(*fullDecisionGraph); + + // 4. Build origToFullDG map + DenseMap origToFullDG; + for (auto [dgBlock, origBlock] : fullDecisionGraph->origMap) + if (origBlock) + origToFullDG[origBlock] = dgBlock; + + // 5. Create cyclic demotion helper + CyclicDemotionHelper demotionHelper( + builder, builder.getContext(), bi, cyclicMgr, origToFullDG, + builder.getUnknownLoc(), insertBlock, pendingMuxOperands, shadow); + + // 6. Build acyclic decision graph + auto acyclicDG = buildDecisionGraph(*locGraph, locConsAcyclicDeps); + + // 7. Pre-register demoted values for high-level variables + SignalRegistry registry; + demotionHelper.preRegisterDemotedValues(acyclicDG, registry); + + // 8. Construct distribution circuit + buildDistributionNetwork(builder, *acyclicDG, bi, registry, + pendingMuxOperands, shadow); + + // 9. CDA on acyclic DG + ControlDependenceAnalysis locCDA(*acyclicDG->region); + DenseSet locConsControlDeps = + locCDA.getAllBlockDeps()[acyclicDG->newCons].allControlDeps; + + // 10. Enumerate paths → expression → circuit + BoolExpression *fBackedge = + enumeratePaths(*acyclicDG, bi, locConsControlDeps); + Value conditionValue = + expressionToCircuit(builder, fBackedge, *acyclicDG, insertBlock, registry, + bi, pendingMuxOperands, shadow); + + // 11. Clean up temporary graphs + acyclicDG->containerOp->erase(); + fullDecisionGraph->containerOp->erase(); + locGraph->containerOp->erase(); + + return conditionValue; +} + +// ===--------------------------------------------------------------------=== // +// LocalCFG and Decision Graph construction +// ===--------------------------------------------------------------------=== // + +/// Reconstructs, as a fresh standalone region, the part of the control-flow +/// graph a token can travel through on its way from `origProd` to `origCons`. +/// +/// Every forward path from producer to consumer is kept; every other edge — +/// paths that miss the consumer, or that return to the producer — is routed to +/// a single `sink` block. "Reaches the consumer" versus "reaches the sink" +/// then captures exactly when the token should be delivered versus discarded. +/// Loops lying on a producer-to-consumer path are kept as real cycles (see the +/// successor handling below); back edges that leave the region are cut to sink. +/// +/// In the result, `newProd` is the entry, `newCons` is the block standing for +/// the consumer, and `sinkBB` is the shared exit. +std::unique_ptr +ftd::buildLocalCFGRegion(OpBuilder &builder, Block *origProd, Block *origCons, + const ftd::BlockIndexing &bi) { + auto L = std::make_unique(); + Location loc = builder.getUnknownLoc(); + + // Setup Region Container + OpBuilder::InsertionGuard guard(builder); + auto funcType = builder.getFunctionType({}, {}); + auto dummyFunc = + builder.create(loc, "__ftd_local_cfg__", funcType); + Region &R = dummyFunc.getBody(); + L->region = &R; + L->containerOp = dummyFunc; + + // Sink Block: The unified exit for all paths (valid or suppressed). + L->sinkBB = new Block(); + R.push_back(L->sinkBB); + L->origMap[L->sinkBB] = nullptr; + + // Producer Block: The entry point of the local CFG. + Block *entry = new Block(); + R.push_back(entry); + L->newProd = entry; + L->origMap[entry] = origProd; + + DenseMap cloned; + DenseSet visited; + // Avoid scheduling the same orig block twice + DenseSet enqueued; + cloned[origProd] = entry; + + // DFS Function + std::function dfs = [&](Block *currOrig, + Block *currNew) { + visited.insert(currOrig); + + auto *term = currOrig->getTerminator(); + + // Dead End: Implicit flow to Sink. + if (!term || term->getNumSuccessors() == 0) { + builder.setInsertionPointToEnd(currNew); + builder.create(loc, L->sinkBB); + return; + } + + // LIST 1: The distinct successors in the NEW Local CFG for the current + // block. Used to construct the BranchOp/CondBranchOp. + SmallVector localSuccessors; + + // LIST 2: The successors that are valid and new, requiring further DFS + // traversal. Stored as pairs: {Original Successor, New Local Block}. + SmallVector, 2> successorsToVisit; + + for (auto it = term->successor_begin(), e = term->successor_end(); it != e; + ++it) { + Block *succOrig = *it; + Block *nextBlockInLocalCFG = + nullptr; // Where the edge points to in the new graph + + // Decide where this edge goes in the new graph. The order of the checks + // matters: the "already cloned" check comes before the back-edge check + // on purpose. An edge to a block we have already cloned reuses that + // clone, which keeps a loop as a real cycle here (later loop handling + // relies on it). Only edges to not-yet-cloned blocks that run backwards + // in block order are treated as region-leaving back edges and cut to + // sink. + + // Edge reaches the consumer: deliver the token here. + if (succOrig == origCons) { + if (succOrig == origProd) { + // Self-loop delivery + if (!L->secondVisitBB) { + L->secondVisitBB = new Block(); + R.push_back(L->secondVisitBB); + L->origMap[L->secondVisitBB] = nullptr; + // Terminate SecondVisit immediately to Sink + OpBuilder::InsertionGuard g(builder); + builder.setInsertionPointToEnd(L->secondVisitBB); + builder.create(loc, L->sinkBB); + } + nextBlockInLocalCFG = L->secondVisitBB; + L->newCons = L->secondVisitBB; + } else { + // Standard delivery + if (!L->newCons || L->newCons == L->secondVisitBB) { + Block *proxy = new Block(); + R.push_back(proxy); + L->origMap[proxy] = succOrig; + L->newCons = proxy; + // Terminate Proxy immediately to Sink + OpBuilder::InsertionGuard g(builder); + builder.setInsertionPointToEnd(proxy); + builder.create(loc, L->sinkBB); + } + nextBlockInLocalCFG = L->newCons; + } + } + // Edge returns to the producer without reaching the consumer: the + // current token would be overwritten by the next one, so discard it. + else if (succOrig == origProd) { + nextBlockInLocalCFG = L->sinkBB; + } + // Edge targets a block we have already cloned: reuse that clone. This is + // what keeps an on-path loop as a cycle in the new graph. If the target + // has not been explored yet, schedule it (once). + else if (cloned.count(succOrig)) { + nextBlockInLocalCFG = cloned[succOrig]; + // If this node hasn't been visited yet, ensure it will be traversed + // once. + if (!visited.count(succOrig) && !enqueued.count(succOrig)) { + enqueued.insert(succOrig); + successorsToVisit.push_back({succOrig, nextBlockInLocalCFG}); + } + } + // Defensive: seen but not cloned (should not happen via this builder). + else if (visited.count(succOrig)) { + nextBlockInLocalCFG = L->sinkBB; + } + // Back edge that leaves the region (targets an earlier, unexplored + // block): the path does not reach the consumer, so cut it to sink. + else if (bi.isLess(succOrig, currOrig)) { + nextBlockInLocalCFG = L->sinkBB; + } + // New forward edge: clone the target and keep exploring. + else { + Block *newSucc = new Block(); + R.push_back(newSucc); + cloned[succOrig] = newSucc; + L->origMap[newSucc] = succOrig; + + nextBlockInLocalCFG = newSucc; + // Schedule this node for DFS visitation (once) + if (!enqueued.count(succOrig)) { + enqueued.insert(succOrig); + successorsToVisit.push_back({succOrig, newSucc}); + } + } + + // Add the determined destination to the list of local successors + localSuccessors.push_back(nextBlockInLocalCFG); + } + + // Create the branch instruction + builder.setInsertionPointToEnd(currNew); + if (localSuccessors.size() == 1) { + builder.create(loc, localSuccessors[0]); + } else if (localSuccessors.size() == 2) { + // Placeholder condition for 2-way branches + Value cond = builder.create(loc, 1, 1); + builder.create(loc, cond, localSuccessors[0], + localSuccessors[1]); + } else { + // Default fall-through for complex control flow + builder.create(loc, L->sinkBB); + } + + // Continue DFS + for (auto &pair : successorsToVisit) { + dfs(pair.first, pair.second); + } + }; + + // Start DFS + dfs(origProd, L->newProd); + + // Finalize Sink + builder.setInsertionPointToEnd(L->sinkBB); + builder.create(loc); + + if (!L->newCons) + L->newCons = L->sinkBB; + + // Compute Topological Order + DenseSet visitedTopo; + SmallVector order; + std::function topo = [&](Block *u) { + if (!u || visitedTopo.contains(u)) + return; + visitedTopo.insert(u); + if (auto *term = u->getTerminator()) + for (auto it = term->successor_begin(), e = term->successor_end(); + it != e; ++it) + topo(*it); + order.push_back(u); + }; + + topo(L->newProd); + std::reverse(order.begin(), order.end()); + L->topoOrder = std::move(order); + + // Physical Reordering + // Reorder blocks in the region list to match the topological order. + // This does not change the graph structure (pointers), only the memory + // layout/print order. + for (Block *b : L->topoOrder) { + if (b != L->sinkBB) { + b->moveBefore(L->sinkBB); + } + } + + return L; +} + +/// Constructs a NEW LocalCFG that represents the Decision Graph. +/// \param rawGraph The source LocalCFG. +/// \param dependencies The set of blocks (from rawGraph) that are relevant +/// decision nodes. +/// \param muxConstraints A map {Block* -> bool} enforcing specific values for +/// blocks. If a block is in this map, the branch corresponding to !value is +/// wired to Sink. +std::unique_ptr +ftd::buildDecisionGraph(const ftd::LocalCFG &rawGraph, + const DenseSet &dependencies, + const DenseMap &muxConstraints) { + + if (!rawGraph.newCons) + return nullptr; + + // 1. NodeSet: Consumer + Sink + Dependencies + DenseSet nodeSet; + nodeSet.insert(rawGraph.newCons); + nodeSet.insert(rawGraph.sinkBB); + for (Block *b : dependencies) { + nodeSet.insert(b); + } + + // 2. Setup New Container + auto newL = std::make_unique(); + OpBuilder builder(rawGraph.containerOp->getContext()); + Location loc = builder.getUnknownLoc(); + + auto funcType = builder.getFunctionType({}, {}); + auto newContainer = + builder.create(loc, "__ftd_decision_graph__", funcType); + Region &newRegion = newContainer.getBody(); + newL->region = &newRegion; + newL->containerOp = newContainer; + + // 3. Create Blocks & Map + // Create a dummy start block that ensures the entry block never has + // back-edges. + Block *dummyStart = new Block(); + newRegion.push_back(dummyStart); + newL->origMap[dummyStart] = nullptr; + + DenseMap oldToNew; + + for (Block *oldBlock : rawGraph.topoOrder) { + if (nodeSet.contains(oldBlock)) { + Block *newBlock = new Block(); + newRegion.push_back(newBlock); + oldToNew[oldBlock] = newBlock; + + if (Block *origIR = rawGraph.origMap.lookup(oldBlock)) { + newL->origMap[newBlock] = origIR; + } else { + newL->origMap[newBlock] = nullptr; + } + + // Set newProd to the first valid block (TopoOrder) + if (newL->newProd == nullptr) { + newL->newProd = newBlock; + } + + if (oldBlock == rawGraph.newCons) + newL->newCons = newBlock; + if (oldBlock == rawGraph.sinkBB) + newL->sinkBB = newBlock; + if (oldBlock == rawGraph.secondVisitBB) + newL->secondVisitBB = newBlock; + } + } + + // Fallback for newProd (unlikely, as newCons/sinkBB are in nodeSet) + if (newL->newProd == nullptr && newRegion.getBlocks().size() > 1) { + newL->newProd = &newRegion.back(); + } + + // 4. Helper: Find Nearest using DFS with Visited Set + auto findNearest = [&](Block *start) -> Block * { + if (!start) + return nullptr; + DenseSet visited; + std::function dfs = [&](Block *curr) -> Block * { + if (!curr) + return nullptr; + if (nodeSet.contains(curr)) + return curr; + if (!visited.insert(curr).second) + return nullptr; // Cycle + for (Block *succ : curr->getSuccessors()) { + if (Block *res = dfs(succ)) + return res; + } + return nullptr; + }; + return dfs(start); + }; + + // 5. Wire the Graph + + // Wire the dummy start unconditionally to the true logic entry + builder.setInsertionPointToEnd(dummyStart); + if (newL->newProd) { + builder.create(loc, newL->newProd); + } else { + builder.create(loc); + } + + for (auto [oldBlock, newBlock] : oldToNew) { + // Sink Logic: Terminate + if (oldBlock == rawGraph.sinkBB) { + builder.setInsertionPointToEnd(newBlock); + builder.create(loc); + continue; + } + + // Consumer Logic: MUST Branch to Sink + if (oldBlock == rawGraph.newCons) { + builder.setInsertionPointToEnd(newBlock); + // In LocalCFG, Consumer always branches to Sink. + // We replicate this connection in the new graph. + Block *newSink = newL->sinkBB; + if (newSink) { + builder.create(loc, newSink); + } else { + // Should not happen if Sink is in nodeSet + builder.create(loc); + } + continue; + } + + // Decision Node Logic + Operation *term = oldBlock->getTerminator(); + if (!term) + continue; + + builder.setInsertionPointToEnd(newBlock); + + if (auto condBr = dyn_cast(term)) { + Block *oldTrue = findNearest(condBr.getTrueDest()); + Block *oldFalse = findNearest(condBr.getFalseDest()); + + Block *newTrue = oldToNew.lookup(oldTrue); + Block *newFalse = oldToNew.lookup(oldFalse); + + // [Safety Wiring] If a path is dead/looping, wire to Sink + if (!newTrue) + newTrue = newL->sinkBB; + if (!newFalse) + newFalse = newL->sinkBB; + + // [Constraint Application] + // If this block has a constraint, wire the invalid path to Sink. + if (muxConstraints.count(oldBlock)) { + bool requiredVal = muxConstraints.lookup(oldBlock); + if (requiredVal) { + // Require True -> Wire False to Sink + newFalse = newL->sinkBB; + } else { + // Require False -> Wire True to Sink + newTrue = newL->sinkBB; + } + } + + Value cond = builder.create(loc, 1, 1); + builder.create(loc, cond, newTrue, newFalse); + } else { + // Non-CondBranch nodes in the decision set (rare) + // Wire to nearest valid successor or Sink + Block *oldTarget = findNearest(term->getSuccessor(0)); + Block *newTarget = oldToNew.lookup(oldTarget); + if (!newTarget) + newTarget = newL->sinkBB; + builder.create(loc, newTarget); + } + } + + // 6. Compute TopoOrder + DenseSet visited; + SmallVector order; + std::function topo = [&](Block *u) { + if (!u || visited.contains(u)) + return; + visited.insert(u); + if (auto *term = u->getTerminator()) + for (auto it = term->successor_begin(); it != term->successor_end(); ++it) + topo(*it); + order.push_back(u); + }; + + if (!newRegion.empty()) { + topo(&newRegion.front()); + std::reverse(order.begin(), order.end()); + newL->topoOrder = std::move(order); + } + + return newL; +} + +// ===--------------------------------------------------------------------=== // +// CyclicDemotionHelper implementation +// ===--------------------------------------------------------------------=== // + +CyclicDemotionHelper::CyclicDemotionHelper( + mlir::OpBuilder &builder, MLIRContext *ctx, const BlockIndexing &bi, + CyclicGraphManager &cyclicMgr, DenseMap &origToFullDG, + Location loc, Block *insertBlock, + DenseMap> *pendingMuxOperands, + ShadowCFG *shadow) + : builder(builder), ctx(ctx), bi(bi), cyclicMgr(cyclicMgr), + origToFullDG(origToFullDG), loc(loc), insertBlock(insertBlock), + pendingMuxOperands(pendingMuxOperands), shadow(shadow) {} + +/// Returns the loop-nesting level at which condition variable `var` is +/// produced. +unsigned CyclicDemotionHelper::getVarNativeLevel(const std::string &var) { + // Map the condition variable to its block, then to the decision-graph + // block, and read that block's loop-nesting level. + auto opt = bi.getBlockFromCondition(var); + if (!opt) + return 0; + auto it = origToFullDG.find(opt.value()); + if (it == origToFullDG.end()) + return 0; + return cyclicMgr.getNestingLevel(it->second); +} + +/// Finds the loop scope at nesting level `level` that contains `origIRBlock`'s +/// decision-graph node, or null if there is none. +LoopScope *CyclicDemotionHelper::findScopeForBlock(Block *origIRBlock, + unsigned level) { + auto it = origToFullDG.find(origIRBlock); + if (it == origToFullDG.end()) + return nullptr; + Block *dgBlock = it->second; + // Walk the scope tree for the scope at `level` that contains the block. + std::function search; + search = [&](ftd::LoopScope *s) -> ftd::LoopScope * { + if (s->level == level && s->allBlocksInclusive.contains(dgBlock)) + return s; + for (auto &sub : s->subLoops) + if (auto *f = search(sub.get())) + return f; + return nullptr; + }; + return search(cyclicMgr.getTopLevelScope()); +} + +/// Demotes `currentValue` (a token produced at loop-nesting level `fromLevel`, +/// anchored at `origBlock`) down one level, so it ends up at the rate of the +/// enclosing level. Builds that level's acyclic CFG, derives the condition that +/// keeps only the token of the iteration on which the loop exits (and on which +/// `origBlock` is reached), and branches the value on it; the kept false result +/// is the demoted value. +Value CyclicDemotionHelper::demoteOneLevel(Value currentValue, Block *origBlock, + unsigned fromLevel) { + ftd::LoopScope *scope = findScopeForBlock(origBlock, fromLevel); + if (!scope) { + llvm::errs() << "[FTD Warning] No LoopScope found at level " << fromLevel + << " for demotion.\n"; + return currentValue; + } + + // 1. Extract acyclic layered CFG for this loop scope + OpBuilder layerBuilder(ctx); + auto levelCFG = cyclicMgr.extractLayeredCFG(scope, layerBuilder); + + // 2. CDA on levelCFG (main suppression: header -> loop exit) + ControlDependenceAnalysis levelCDA(*levelCFG->region); + DenseSet levelDeps = + levelCDA.getAllBlockDeps()[levelCFG->newCons].allControlDeps; + + // 3. Create level-local registry and pre-register all dep variables + // at their correct level (demoting higher-level ones recursively) + SignalRegistry levelRegistry; + for (Block *depBlock : levelDeps) { + Block *depOrigBlock = levelCFG->origMap.lookup(depBlock); + if (!depOrigBlock) + continue; + std::string depVar = bi.getBlockCondition(depOrigBlock); + if (depVar.empty()) + continue; + + unsigned depNative = getVarNativeLevel(depVar); + Value depValue; + if (depNative <= fromLevel) { + // Variable is at this level or outer — use original value + depValue = + getOriginalValue(builder, depVar, bi, pendingMuxOperands, shadow); + } else { + // Variable is from a deeper loop — demote to this level + depValue = getValueAtLevel(depVar, fromLevel); + } + if (depValue) { + if (!depValue.getType().isa()) + depValue.setType(ftd::channelifyType(depValue.getType())); + levelRegistry.registerSignal(depVar, {}, depValue); + } + } + + // 4. Distribution on levelCFG + buildDistributionNetwork(builder, *levelCFG, bi, levelRegistry, + pendingMuxOperands, shadow); + + // 5. Main suppression expression: header -> loop exit + BoolExpression *fCons = enumeratePaths(*levelCFG, bi, levelDeps); + fCons = fCons->boolMinimize(); + BoolExpression *fSup = fCons->boolNegate()->boolMinimize(); + + // 6. DP suppression: header -> value's block + BoolExpression *fSupDP = BoolExpression::boolZero(); + DenseMap rankDP; // pre-computed rank for DP graph + Block *origHeader = levelCFG->origMap.lookup(levelCFG->newProd); + + if (origHeader && origHeader != origBlock) { + OpBuilder dpBuilder(ctx); + auto dpLocGraph = buildLocalCFGRegion(dpBuilder, origHeader, origBlock, bi); + if (dpLocGraph->newCons) { + ControlDependenceAnalysis dpCDA(*dpLocGraph->region); + auto dpDepsTmp = + dpCDA.getAllBlockDeps()[dpLocGraph->newCons].allControlDeps; + auto dpDG = buildDecisionGraph(*dpLocGraph, dpDepsTmp); + ControlDependenceAnalysis dpDGCDA(*dpDG->region); + auto dpDeps = dpDGCDA.getAllBlockDeps()[dpDG->newCons].allControlDeps; + BoolExpression *fConsDP = enumeratePaths(*dpDG, bi, dpDeps); + fSupDP = fConsDP->boolMinimize()->boolNegate()->boolMinimize(); + + // Ensure all DP expression variables are in the level registry + std::set dpVars = fSupDP->getVariables(); + for (const auto &v : dpVars) { + if (!levelRegistry.lookup(v, {})) { + unsigned vNative = getVarNativeLevel(v); + Value vVal; + if (vNative <= fromLevel) + vVal = getOriginalValue(builder, v, bi, pendingMuxOperands, shadow); + else + vVal = getValueAtLevel(v, fromLevel); + if (vVal) { + if (!vVal.getType().isa()) + vVal.setType(ftd::channelifyType(vVal.getType())); + levelRegistry.registerSignal(v, {}, vVal); + } + } + } + + // Pre-compute DP rank before erasing + rankDP = computeTopoRank(*dpDG); + + dpDG->containerOp->erase(); + } + dpLocGraph->containerOp->erase(); + } + + // 7. Build suppression circuit + Value result = currentValue; + if (!currentValue.getType().isa()) + currentValue.setType(ftd::channelifyType(currentValue.getType())); + + if (fSup->type != experimental::boolean::ExpressionType::Zero) { + DenseMap rank = computeTopoRank(*levelCFG); + Value branchCond = + expressionToCircuit(builder, fSup, rank, insertBlock, levelRegistry, bi, + pendingMuxOperands, shadow); + + // Cascaded DP filter + if (fSupDP->type != experimental::boolean::ExpressionType::Zero) { + Value dpCond = + expressionToCircuit(builder, fSupDP, rankDP, insertBlock, + levelRegistry, bi, pendingMuxOperands, shadow); + auto dpBranch = builder.create( + loc, ftd::getListTypes(branchCond.getType()), dpCond, branchCond); + dpBranch->setAttr(FTD_OP_TO_SKIP, builder.getUnitAttr()); + setBBAttr(dpBranch, insertBlock, builder); + branchCond = dpBranch.getFalseResult(); + } + + auto branchOp = builder.create( + loc, ftd::getListTypes(currentValue.getType()), branchCond, + currentValue); + branchOp->setAttr(FTD_OP_TO_SKIP, builder.getUnitAttr()); + setBBAttr(branchOp, insertBlock, builder); + result = branchOp.getFalseResult(); + } + + levelCFG->containerOp->erase(); + return result; +} + +/// Returns the wire carrying condition variable `varName` valid at loop-nesting +/// level `targetLevel`, building and caching it on demand. If the variable is +/// already produced at or below `targetLevel`, its original wire is used as is; +/// otherwise the value is taken one level up and demoted down a single level, +/// recursing until the variable's native level is reached. Results are +/// memorized in demotionCache, so each (variable, level) wire is built once and +/// shared. +Value CyclicDemotionHelper::getValueAtLevel(const std::string &varName, + unsigned targetLevel) { + auto key = std::make_pair(varName, targetLevel); + auto it = demotionCache.find(key); + if (it != demotionCache.end()) + return it->second; + + unsigned native = getVarNativeLevel(varName); + Value val; + if (native <= targetLevel) { + // Variable is at or below target level — use original value + val = getOriginalValue(builder, varName, bi, pendingMuxOperands, shadow); + if (val && !val.getType().isa()) + val.setType(ftd::channelifyType(val.getType())); + } else { + // Recursively get value one level above, then demote + Value higher = getValueAtLevel(varName, targetLevel + 1); + auto blockOpt = bi.getBlockFromCondition(varName); + assert(blockOpt.has_value() && "Variable block not found"); + val = demoteOneLevel(higher, blockOpt.value(), targetLevel + 1); + } + + demotionCache[key] = val; + return val; +} + +/// Computes the suppression condition for a value flowing from +/// `producerBlock` to `consumerBlock` and returns the basic blocks of the +/// branch conditions it involves. +static SmallVector +collectSuppressionConditionBlocks(Block *producerBlock, Block *consumerBlock, + const ftd::BlockIndexing &bi) { + SmallVector condBlocks; + if (!producerBlock || !consumerBlock || producerBlock == consumerBlock) + return condBlocks; + + OpBuilder tmp(producerBlock->getParent()->getContext()); + auto loc = ftd::buildLocalCFGRegion(tmp, producerBlock, consumerBlock, bi); + if (loc->newCons && loc->newCons != loc->sinkBB) { + ControlDependenceAnalysis locCDA(*loc->region); + DenseSet deps = + locCDA.getAllBlockDeps()[loc->newCons].allControlDeps; + auto dg = ftd::buildDecisionGraph(*loc, deps); + + ControlDependenceAnalysis dgCDA(*dg->region); + DenseSet dgDeps = + dgCDA.getAllBlockDeps()[dg->newCons].allControlDeps; + + BoolExpression *fCons = enumeratePaths(*dg, bi, dgDeps); + BoolExpression *fSup = fCons->boolMinimize()->boolNegate()->boolMinimize(); + + for (const std::string &var : fSup->getVariables()) + if (auto blkOpt = bi.getBlockFromCondition(var)) + condBlocks.push_back(blkOpt.value()); + + dg->containerOp->erase(); + } + loc->containerOp->erase(); + return condBlocks; +} + +// ===--------------------------------------------------------------------=== // +// insertDirectSuppression — main suppression entry point +// ===--------------------------------------------------------------------=== // + +/// Builds the suppression circuit for a single producer value consumed by +/// `consumer`. In outline: pick the block the analysis starts from (for a gamma +/// consumer, the condition that effectively controls the delivery); reconstruct +/// the control flow from there down to the consumer; derive the condition under +/// which the token must be discarded (for a gamma, the token is kept only on +/// paths that select this input); then emit the mux-tree circuit for the +/// condition and branch the producer token on it so the unused token is +/// dropped. +void ftd::insertDirectSuppression(mlir::OpBuilder &builder, + handshake::FuncOp &funcOp, + Operation *consumer, Value connection, + ftd::ShadowCFG &shadow) { + Region &shadowRegion = shadow.getRegion(); + ftd::BlockIndexing bi(shadowRegion); + + // Map producer and consumer to shadow blocks via handshake.bb + unsigned prodBBIdx = 0; + if (auto *defOp = connection.getDefiningOp()) + if (auto attr = defOp->getAttrOfType("handshake.bb")) + prodBBIdx = attr.getUInt(); + unsigned consBBIdx = 0; + if (auto attr = consumer->getAttrOfType("handshake.bb")) + consBBIdx = attr.getUInt(); + IntegerAttr prodBBAttr = ftd::getBBIndexAttr(builder.getContext(), prodBBIdx); + // Suppression ops are tagged to the producer's block by default; if the + // consumer is an enclosing loop header, place them at the loop exit instead. + IntegerAttr targetBBAttr = prodBBAttr; + Block *producerIRBlock = connection.getDefiningOp() + ? connection.getDefiningOp()->getBlock() + : consumer->getBlock(); + + Block *entryBlock = shadow.getBlock(0); + Block *producerBlock = shadow.getBlock(prodBBIdx); + Block *consumerBlock = shadow.getBlock(consBBIdx); + if (IntegerAttr loopExitBBAttr = getFirstLoopExitBBAttrIfHeaderConsumer( + builder, shadowRegion, producerBlock, consumerBlock, bi, shadow)) + targetBBAttr = loopExitBBAttr; + // The block the suppression analysis starts from. It dominates the producer + // and every condition block the suppression will depend on. + Block *dominatorBlock = producerBlock; + + // Account for the condition of a Mux only if it corresponds to a GAMMA GSA + // gate + bool deliverToGamma = llvm::isa(consumer) && + consumer->hasAttr(FTD_EXPLICIT_GAMMA) && + (producerBlock != consumerBlock || + connection.getDefiningOp()->hasAttr(FTD_EXPLICIT_MU)); + + // If producer is unreachable, the suppression is not needed. + if (!isReachable(entryBlock, producerBlock)) { + return; + } + + // Reads an operation's basic-block index from its handshake.bb attribute. + auto getBB = [](Operation *op) -> unsigned { + auto attr = op->getAttrOfType("handshake.bb"); + return attr ? attr.getUInt() : 0; + }; + + // When the consumer is a gamma mux tree, the producer feeds only one input + // on it. Walk the tree to find the highest branch that can still route to + // the producer block on both of its sides, which indicates it effectively + // controls the delivery. + if (deliverToGamma) { + Operation *currentMuxOp = consumer; + Value currentConnection = connection; + Block *lastValidDominator = producerBlock; + // Acyclic reachability check: only follows forward edges (skips back-edges) + // using the block ordering from BlockIndexing. + auto isReachableAcyclic = [&](Block *start, Block *end) -> bool { + if (start == end) + return true; + DenseSet visited; + SmallVector stack; + stack.push_back(start); + visited.insert(start); + while (!stack.empty()) { + Block *curr = stack.pop_back_val(); + for (Block *succ : curr->getSuccessors()) { + if (bi.isLess(succ, curr) || succ == curr) + continue; + if (succ == end) + return true; + if (visited.insert(succ).second) + stack.push_back(succ); + } + } + return false; + }; + while (true) { + // The value reaches this mux either as its select (operand 0) or as a + // data input (operand 1/2). If it is a data input, this mux's condition + // must be taken into account; if it is the select, there is nothing to + // take into account and we move on to the next mux. + if (currentMuxOp->getOperand(0) != currentConnection) { + // Connection is a data input — check this Mux's condition block + Value condition = currentMuxOp->getOperand(0); + Block *condBlock = returnMuxConditionBlock(condition, shadow); + + // condBlock is a candidate only if the producer is reachable + // from both of its successors (following forward edges only). + bool bothReach = + condBlock && condBlock->getNumSuccessors() >= 2 && + isReachableAcyclic(condBlock->getSuccessor(0), producerBlock) && + isReachableAcyclic(condBlock->getSuccessor(1), producerBlock); + + // Of the qualifying condition blocks, keep the one earliest in block + // order: it dominates the others and reaches the producer on both of + // its branches. + if (bothReach && bi.isLess(condBlock, lastValidDominator)) { + lastValidDominator = condBlock; + } + } + + // Move on to the next mux of the same gamma tree (its muxes share one + // basic-block) that this mux's output feeds into. + Operation *nextMuxOp = nullptr; + Value currentResult = currentMuxOp->getResult(0); + for (auto *user : currentResult.getUsers()) { + if (llvm::isa(user) && + user->hasAttr(FTD_EXPLICIT_GAMMA) && + getBB(user) == getBB(currentMuxOp)) { + unsigned cnt = 0; + if (user->getOperand(1) == currentResult) + cnt++; + if (user->getOperand(2) == currentResult) + cnt++; + if (cnt != 2) { + nextMuxOp = user; + break; + } + } + } + + if (!nextMuxOp) + break; + currentConnection = currentMuxOp->getResult(0); + currentMuxOp = nextMuxOp; + } + dominatorBlock = lastValidDominator; + + // dominatorBlock must dominate the producer and every block whose branch + // condition the suppression depends on, so take the nearest common + // dominator of all of them. + SmallVector exprCondBlocks = + collectSuppressionConditionBlocks(producerBlock, consumerBlock, bi); + if (!exprCondBlocks.empty()) { + DominanceInfo domInfo; + Block *bd = + domInfo.findNearestCommonDominator(producerBlock, dominatorBlock); + for (Block *cb : exprCondBlocks) { + if (!bd) + break; + bd = domInfo.findNearestCommonDominator(bd, cb); + } + if (bd) + dominatorBlock = bd; + } + } + + // buildLocalCFGRegion builds a throwaway region used only for analysis; give + // it its own builder so those temporary operations stay separate from the IR + // we are actually editing (they are erased manually later). + OpBuilder tmpBuilder(funcOp.getContext()); + // The local CFG only captures the part between the dominator block (functions + // as a producer block) and the consumer block. + auto locGraph = + buildLocalCFGRegion(tmpBuilder, dominatorBlock, consumerBlock, bi); + + ControlDependenceAnalysis locCDA(*locGraph->region); + // The condition blocks the consumer's reachability depends on, from the + // dominator block; the suppression condition is built over exactly these. + // Comprised of the consumer's own control dependences, and (for a gamma) + // extended below with the condition blocks driving the gamma's selects and + // their dependences. + DenseSet locConsControlDepsFull = + locCDA.getAllBlockDeps()[locGraph->newCons].allControlDeps; + // The condition blocks driving the gamma's selects along the consumer input's + // path — the conditions under which the consumer input is the one selected. + DenseSet muxConditionSet; + + // For each such condition block, the value its condition must take for the + // consumer input to be selected. Used below to prune the decision graph so + // the token is kept only on paths that actually route to the producer. + DenseMap muxConstraints; + + // Walk the gamma tree again, now that the dominator block and local CFG are + // known, recording for every mux the producer flows through the condition + // driving that mux and the value it must take to select the input to be + // consumed. This fills muxConditionSet and muxConstraints and extends + // locConsControlDepsFull, capturing that the producer is consumed only under + // this combination of selects. + if (deliverToGamma) { + Operation *currentMuxOp = consumer; + Value currentConnection = connection; + bool isChainActive = true; + + while (isChainActive) { + bool isDataInput = false; + bool requiredVal = true; + + // Check how the connection enters the current Mux + // If operand(0) == currentConnection, it is the condition input. + // In that case, isDataInput remains false, and we simply traverse + // to the output to find the next mux in the chain. + if (!(currentMuxOp->getOperand(0) == currentConnection)) { + if (currentMuxOp->getOperand(1) == currentConnection) { + // Input 1 is the FALSE input + isDataInput = true; + requiredVal = false; + } else if (currentMuxOp->getOperand(2) == currentConnection) { + // Input 2 is the TRUE input + isDataInput = true; + requiredVal = true; + } + } + + if (isDataInput) { + // 1. Get the condition value driving this Mux + Value muxCondition = currentMuxOp->getOperand(0); + + // Trace back through suppression branches + // to find the original source of the condition signal. + while (Operation *defOp = muxCondition.getDefiningOp()) { + if (llvm::isa(defOp) && + defOp->hasAttr(FTD_OP_TO_SKIP)) { + // Move up to the data input of the branch + muxCondition = defOp->getOperand(1); + } else { + break; + } + } + + // 2. Identify the Original Block defining this condition variable + // (Using the provided helper function) + Block *muxConditionBlock = + returnMuxConditionBlock(muxCondition, shadow); + muxConditionSet.insert(muxConditionBlock); + + // 3. Find the corresponding Block in the Local CFG + // Since locGraph->origMap maps Local->Original, we iterate to reverse + // lookup. + Block *condBlockLocal = nullptr; + for (auto it : locGraph->origMap) { + if (it.second == muxConditionBlock) { + condBlockLocal = it.first; + break; + } + } + + // 4. Add to dependencies and record requirement + if (condBlockLocal) { + if (bi.isLess(muxConditionBlock, dominatorBlock)) + continue; + // Add this block and its all-dependency blocks to the dependency set + // so path enumeration observes it. + locConsControlDepsFull.insert(condBlockLocal); + for (Block *dep : + locCDA.getAllBlockDeps()[condBlockLocal].allControlDeps) { + locConsControlDepsFull.insert(dep); + } + + // Record the specific value required (True/False) to pass this Mux + muxConstraints[condBlockLocal] = requiredVal; + } + } + + // 5. Traverse Downstream (Search for cascaded Gamma Muxes) + Operation *nextMuxOp = nullptr; + Value currentResult = currentMuxOp->getResult(0); + + for (auto *user : currentResult.getUsers()) { + // Check if user is a Mux in the same block marked as Gamma + if (llvm::isa(user) && + user->hasAttr(FTD_EXPLICIT_GAMMA) && + getBB(user) == getBB(currentMuxOp)) { + // Count how many data inputs use currentResult + unsigned connectionCount = 0; + if (user->getOperand(1) == currentResult) + connectionCount++; + if (user->getOperand(2) == currentResult) + connectionCount++; + + // We only proceed if at most ONE data input comes from the previous + // mux. Otherwise, the user is a temporary MUX which we don't care + // about. + if (connectionCount != 2) { + nextMuxOp = user; + // Update connection for next iteration + currentConnection = currentResult; + break; + } + } + } + + if (nextMuxOp) { + currentMuxOp = nextMuxOp; + } else { + isChainActive = false; + } + } + } + + // Early exit: no condition affects this delivery, so there is nothing to + // build a suppression from. Either the consumer is never reached on a kept + // path (discard the token unconditionally, below) or it is always reached + // (nothing to do). + if (locConsControlDepsFull.empty()) { + // The consumer is never reached on any kept path: discard the token on + // every execution by holding the suppress branch's select permanently true. + if (locGraph->newCons == locGraph->sinkBB) { + builder.setInsertionPointToStart(consumer->getBlock()); + auto src = builder.create(consumer->getLoc()); + src->setAttr(FTD_OP_TO_SKIP, builder.getUnitAttr()); + src->setAttr("handshake.bb", targetBBAttr); + auto cstAttr = builder.getIntegerAttr(builder.getIntegerType(1), 1); + auto constOp = builder.create( + consumer->getLoc(), cstAttr, src.getResult()); + constOp->setAttr(FTD_OP_TO_SKIP, builder.getUnitAttr()); + constOp->setAttr("handshake.bb", targetBBAttr); + + Value supData = connection; + auto branchOp = builder.create( + consumer->getLoc(), ftd::getListTypes(supData.getType()), + constOp.getResult(), supData); + branchOp->setAttr(FTD_OP_TO_SKIP, builder.getUnitAttr()); + branchOp->setAttr("handshake.bb", targetBBAttr); + + for (auto &use : llvm::make_early_inc_range(connection.getUses())) { + if (use.getOwner() != consumer) + continue; + use.set(branchOp.getFalseResult()); + } + locGraph->containerOp->erase(); + return; + } + locGraph->containerOp->erase(); + return; + } + + // From here we build the actual suppression: keep the producer's token only + // when the consumer would really use it, and discard it otherwise. (For a + // gamma consumer the full set of conditions routes the select signals, while + // muxConstraints narrows the kept-token condition down to this specific + // input; for a non-gamma consumer muxConstraints is empty and the two + // coincide.) + // + // registry caches the condition signals we materialize, so each is built once + // and reused. The decision graph (built just below) is a view of when + // the consumer is reached, derived from the local CFG and the condition + // blocks it depends on. + SignalRegistry registry; + builder.setInsertionPointToStart(consumer->getBlock()); + auto fullDecisionGraph = + buildDecisionGraph(*locGraph, locConsControlDepsFull); + + // The decision graph may contain loops; this discovers their nesting so the + // graph can later be peeled into acyclic, per-loop-level views. + ftd::CyclicGraphManager cyclicMgr(*fullDecisionGraph); + + // Lookup from an original IR block to its node in the decision graph. + DenseMap origToFullDG; + for (auto [dgBlock, origBlock] : fullDecisionGraph->origMap) + if (origBlock) + origToFullDG[origBlock] = dgBlock; + + // Create the cyclic demotion helper + CyclicDemotionHelper demotionHelper( + builder, funcOp.getContext(), bi, cyclicMgr, origToFullDG, + consumer->getLoc(), consumer->getBlock(), nullptr, &shadow); + + // Level 0: Extract acyclic layered CFG from the full decision graph. + // This cuts all back-edges, giving a DAG where CDA produces correct + // forward dependencies. + OpBuilder level0Builder(funcOp.getContext()); + auto level0CFG = + cyclicMgr.extractLayeredCFG(cyclicMgr.getTopLevelScope(), level0Builder); + + // CDA on the acyclic level 0 CFG — allControlDeps == forwardControlDeps + // because the graph is a DAG. + ControlDependenceAnalysis level0CDA(*level0CFG->region); + DenseSet level0Deps = + level0CDA.getAllBlockDeps()[level0CFG->newCons].allControlDeps; + + // Translate mux condition blocks (original IR) into level0CFG and insert + // them into the dependency set. + for (Block *muxCondOrigIR : muxConditionSet) { + for (auto &[l0Block, l0Orig] : level0CFG->origMap) { + if (l0Orig == muxCondOrigIR) { + level0Deps.insert(l0Block); + for (Block *dep : level0CDA.getAllBlockDeps()[l0Block].allControlDeps) + level0Deps.insert(dep); + break; + } + } + } + + // Translate mux constraints from locGraph blocks to level0CFG blocks. + DenseMap level0MuxConstraints; + for (auto &[locBlock, reqVal] : muxConstraints) { + Block *origIR = locGraph->origMap.lookup(locBlock); + if (!origIR) + continue; + for (auto &[l0Block, l0Orig] : level0CFG->origMap) { + if (l0Orig == origIR) { + level0MuxConstraints[l0Block] = reqVal; + break; + } + } + } + + // The decision graph at loop-nesting level 0 (all back edges cut), left + // unconstrained. The distribution logic must run on this full graph to + // keep the right semantics; the mux-constraint trimming is applied only + // afterwards, on the separate constrained graph. + auto level0FullDG = buildDecisionGraph(*level0CFG, level0Deps); + + // Variables produced inside loops live at a deeper loop-nesting level; demote + // them down to level 0 ahead of time and cache them, so the routing and the + // expression below consume the level-0 wires. + demotionHelper.preRegisterDemotedValues(level0FullDG, registry); + + // Build distribution on level 0 + buildDistributionNetwork(builder, *level0FullDG, bi, registry, nullptr, + &shadow); + + // The same level-0 graph, now pruned by muxConstraints so it keeps only the + // paths that actually select the input to be consumed. The suppression + // condition is read from this constrained graph. + auto level0ConstrainedDG = + buildDecisionGraph(*level0CFG, level0Deps, level0MuxConstraints); + + ControlDependenceAnalysis decCDA(*level0ConstrainedDG->region); + DenseSet constrainedDeps = + decCDA.getAllBlockDeps()[level0ConstrainedDG->newCons].allControlDeps; + + // fCons is the condition under which the consumer actually uses the token; + // its negation, fSup, is when the token must be discarded. + BoolExpression *fCons = + enumeratePaths(*level0ConstrainedDG, bi, constrainedDeps); + + fCons = fCons->boolMinimize(); + BoolExpression *fSup = fCons->boolNegate(); + fSup = fSup->boolMinimize(); + + // Build the circuit that computes, for a `start`->`target` stretch, the + // expression under which the token must be discarded because it does not + // reach `target`, and return that expression wire. Returns null when the + // token is never discarded. + // + // The result is always reduced to loop-nesting level 0: + // - The start->target decision graph may contain loops, so its back edges + // are + // cut (extractLayeredCFG on the top scope) to get an acyclic level-0 + // graph. + // - Every condition signal it consumes is taken at level 0: the demoted + // level-0 wires are pulled through the shared demotionHelper, so a + // condition already demoted elsewhere is reused (the cache returns the + // same wire, forked here) and no deeper-level token leaks in. + // - Distribution is rebuilt into a separate registry and cannot be + // shared; only the demoted condition wire is shared. + auto buildLevel0SuppressionSelect = [&](Block *start, + Block *target) -> Value { + if (!start || !target || start == target) + return Value(); + + OpBuilder cfgBuilder(funcOp.getContext()); + auto locG = buildLocalCFGRegion(cfgBuilder, start, target, bi); + if (!locG->newCons) { + locG->containerOp->erase(); + return Value(); + } + + // Cyclic decision graph, then its level-0 (acyclic) reduction. + ControlDependenceAnalysis locCDAlocal(*locG->region); + DenseSet depsTmp = + locCDAlocal.getAllBlockDeps()[locG->newCons].allControlDeps; + auto fullDG = buildDecisionGraph(*locG, depsTmp); + + CyclicGraphManager cyc(*fullDG); + OpBuilder l0Builder(funcOp.getContext()); + auto level0CFGlocal = + cyc.extractLayeredCFG(cyc.getTopLevelScope(), l0Builder); + + ControlDependenceAnalysis l0CDA(*level0CFGlocal->region); + DenseSet l0Deps = + l0CDA.getAllBlockDeps()[level0CFGlocal->newCons].allControlDeps; + auto level0DG = buildDecisionGraph(*level0CFGlocal, l0Deps); + + ControlDependenceAnalysis dgCDA(*level0DG->region); + DenseSet dgDeps = + dgCDA.getAllBlockDeps()[level0DG->newCons].allControlDeps; + + // Private registry: reuse the already-demoted level-0 condition wires + // (shared via the demotion cache in demotionHelper), then build this + // delivery's own distribution branches into it. + SignalRegistry localRegistry; + demotionHelper.preRegisterDemotedValues(level0DG, localRegistry); + buildDistributionNetwork(builder, *level0DG, bi, localRegistry, nullptr, + &shadow); + + BoolExpression *fConsLocal = enumeratePaths(*level0DG, bi, dgDeps); + BoolExpression *fSupLocal = + fConsLocal->boolMinimize()->boolNegate()->boolMinimize(); + + Value result; + if (fSupLocal->type != experimental::boolean::ExpressionType::Zero) { + DenseMap rk = computeTopoRank(*level0DG); + result = expressionToCircuit(builder, fSupLocal, rk, producerIRBlock, + localRegistry, bi, nullptr, &shadow, + targetBBAttr); + } + + level0DG->containerOp->erase(); + level0CFGlocal->containerOp->erase(); + fullDG->containerOp->erase(); + locG->containerOp->erase(); + return result; + }; + + // The producer-reaching filter. The dominator can sit above the producer, so + // the producer does not fire on every path that reaches the dominator; this + // computes the condition for exactly those dominator iterations on which the + // producer never fires, so the main condition can be trimmed by it. It is + // built at level 0 and reuses the already-demoted condition wires, so it + // injects no deeper-level tokens into the cascade below. Null when the + // dominator is the producer, or the producer is always reached (the filter is + // then optimized away). + Value dpBranchCond; + if (dominatorBlock != producerBlock) + dpBranchCond = buildLevel0SuppressionSelect(dominatorBlock, producerBlock); + + Value supData = connection; + + level0CFG->containerOp->erase(); + level0FullDG->containerOp->erase(); + level0ConstrainedDG->containerOp->erase(); + + // Loop filter + value demotion. + // + // When the producer sits in a deeper loop than the dominator block, it fires + // several times per delivery and produces extra tokens that the main + // suppression branch below cannot match. This is resolved with the demotion + // mechanism applied to the producer's value. + // + // We decide whether the producer is in a loop deeper than the dominator. If + // it is, two cases arise depending on whether the producer block appears as a + // node of the decision graph built for this dominator->consumer delivery: + // + // (A) Producer IS a decision-graph node. The producer value is demoted level + // by level using that node's demotion logic, and the demoted value then + // enters the normal suppression branch below. + // + // (B) Producer is NOT a decision-graph node. It may lie on + // deeper simple loops that are post-dominated by some decision-graph + // node. We first run a simple producer->postdominator suppression (which + // never involves a gamma tree or a cyclic decision graph — either simple + // loops or a straight pass-through whose condition is constant true) to + // filter the simple loops, reusing the ordinary suppression flow. The + // filtered value is then demoted using that postdominator node's + // demotion logic. + if (producerBlock && dominatorBlock) { + DominanceInfo lfDom; + mlir::CFGLoopInfo lfLoopInfo(lfDom.getDomTree(&shadowRegion)); + auto loopDepth = [&](Block *b) -> unsigned { + if (mlir::CFGLoop *l = lfLoopInfo.getLoopFor(b)) + return l->getLoopDepth(); + return 0; + }; + + unsigned domLevel = loopDepth(dominatorBlock); + unsigned prodLevel = loopDepth(producerBlock); + + if (prodLevel > domLevel) { + // Original IR block whose decision-graph node drives the demotion + // cascade, and the block the demotion is anchored on. + Block *demoteBlock = nullptr; + + if (origToFullDG.count(producerBlock)) { + // Case (A): the producer is itself a decision-graph node. + demoteBlock = producerBlock; + } else { + // Case (B): the producer is not a decision-graph node. Find the nearest + // node that post-dominates it. + mlir::PostDominanceInfo lfPostDom; + auto &pdt = lfPostDom.getDomTree(&shadowRegion); + Block *pdNode = nullptr; + if (auto *node = pdt.getNode(producerBlock)) { + // Climb the post-dominator tree to the nearest block that is a + // decision-graph node. + for (node = node->getIDom(); node && node->getBlock(); + node = node->getIDom()) { + if (origToFullDG.count(node->getBlock())) { + // accept only decision-graph nodes + pdNode = node->getBlock(); + break; + } + } + } + + if (pdNode && pdNode != producerBlock) { + // Simple suppression from the producer to the postdominator node. + // This never involves a gamma tree or a cyclic decision graph, so the + // plain enumeratePaths -> expressionToCircuit flow is sufficient. + // When the producer reaches the node unconditionally, lfSup is + // constant zero and the filter is optimized away. + OpBuilder lfBuilder(funcOp.getContext()); + auto lfLoc = + buildLocalCFGRegion(lfBuilder, producerBlock, pdNode, bi); + if (lfLoc->newCons) { + ControlDependenceAnalysis lfLocCDA(*lfLoc->region); + DenseSet lfDepsTmp = + lfLocCDA.getAllBlockDeps()[lfLoc->newCons].allControlDeps; + auto lfDG = buildDecisionGraph(*lfLoc, lfDepsTmp); + + ControlDependenceAnalysis lfDGCDA(*lfDG->region); + DenseSet lfDeps = + lfDGCDA.getAllBlockDeps()[lfDG->newCons].allControlDeps; + + BoolExpression *lfCons = enumeratePaths(*lfDG, bi, lfDeps); + BoolExpression *lfSup = + lfCons->boolMinimize()->boolNegate()->boolMinimize(); + + if (lfSup->type != experimental::boolean::ExpressionType::Zero) { + DenseMap lfRank = computeTopoRank(*lfDG); + // The filter circuit belongs to the producer's basic block. + Value lfCond = expressionToCircuit(builder, lfSup, lfRank, + producerIRBlock, registry, bi, + nullptr, &shadow, prodBBAttr); + auto lfBranch = builder.create( + consumer->getLoc(), ftd::getListTypes(supData.getType()), + lfCond, supData); + lfBranch->setAttr(FTD_OP_TO_SKIP, builder.getUnitAttr()); + lfBranch->setAttr("handshake.bb", prodBBAttr); + // Pass the token through only when the node is reached. + supData = lfBranch.getFalseResult(); + } + lfDG->containerOp->erase(); + } + lfLoc->containerOp->erase(); + + // The filtered value is demoted from the postdominator node onward. + demoteBlock = pdNode; + } + } + + // Demote the (possibly simple-filtered) producer value level by level, + // using the decision-graph nesting of demoteBlock. demoteOneLevel keeps, + // at each level, only the token of the iteration on which that level's + // loop exits (and on which demoteBlock is reached), so after the cascade + // the value is valid at the dominator-to-consumer level and pairs + // one-to-one with the main suppression branch below. + if (demoteBlock && origToFullDG.count(demoteBlock)) { + unsigned demoteLevel = + cyclicMgr.getNestingLevel(origToFullDG[demoteBlock]); + for (unsigned lvl = demoteLevel; lvl >= 1; --lvl) + supData = demotionHelper.demoteOneLevel(supData, demoteBlock, lvl); + } + } + } + + // The decision graph (and the cyclic machinery built on top of it) had to + // stay alive for the demotion cascade above; erase it now. + fullDecisionGraph->containerOp->erase(); + + // fSup is the condition under which the consumer does not use the token. When + // it can be non-zero, build its circuit (branchCond) and branch the producer + // token on it, dropping the token on exactly the executions where the + // consumer would not use it. (When fSup is constant zero the consumer always + // uses the token, so no suppress branch is emitted.) + if (fSup->type != experimental::boolean::ExpressionType::Zero) { + IntegerAttr connectionBBAttr = targetBBAttr; + DenseMap rank = computeTopoRank(*locGraph); + Value branchCond = + expressionToCircuit(builder, fSup, rank, producerIRBlock, registry, bi, + nullptr, &shadow, connectionBBAttr); + + // The dominator may sit above the producer, so branchCond (fSup) carries + // one token per dominator activation, while the producer fires only on some + // of them. dpBranchCond (fDP) is true exactly on the activations where the + // producer is not reached; branching branchCond on it and keeping the false + // result drops those activations, leaving branchCond paired one-to-one with + // the producer tokens before it drives the suppress branch below. + if (dpBranchCond) { + auto dpBranchOp = builder.create( + consumer->getLoc(), ftd::getListTypes(branchCond.getType()), + dpBranchCond, branchCond); + dpBranchOp->setAttr(FTD_OP_TO_SKIP, builder.getUnitAttr()); + dpBranchOp->setAttr("handshake.bb", targetBBAttr); + branchCond = dpBranchOp.getFalseResult(); + } + + // Suppress branch: discard the producer token whenever branchCond holds; + // the consumer keeps it only via the false result. + auto branchOp = builder.create( + consumer->getLoc(), ftd::getListTypes(supData.getType()), branchCond, + supData); + branchOp->setAttr(FTD_OP_TO_SKIP, builder.getUnitAttr()); + branchOp->setAttr("handshake.bb", targetBBAttr); + supData = branchOp.getFalseResult(); + + // Take into account the possibility of a mux to get the condition input + // also as data input. In this case, the data input can be optimized to + // a constant value, since it is always selected only when its own value + // is true or false. + for (auto &use : llvm::make_early_inc_range(connection.getUses())) { + if (use.getOwner() != consumer) + continue; + if (llvm::isa(consumer) && + consumer->getOperand(0) == connection && + use.getOperandNumber() != 0) { + auto src = builder.create(consumer->getLoc()); + setBBAttrWithFallback(src, targetBBAttr, producerIRBlock, builder); + auto innerType = + connection.getType().cast().getDataType(); + auto attr = + builder.getIntegerAttr(innerType, (use.getOperandNumber() == 2)); + auto cst = builder.create( + consumer->getLoc(), connection.getType(), attr, src.getResult()); + cst->setAttr(FTD_OP_TO_SKIP, builder.getUnitAttr()); + setBBAttrWithFallback(cst, targetBBAttr, producerIRBlock, builder); + use.set(cst.getResult()); + continue; + } + use.set(supData); + } + } else if (supData != connection) { + // No main suppression was needed, but the loop filter rewrote the + // producer token; route the filtered token to the consumer. + for (auto &use : llvm::make_early_inc_range(connection.getUses())) { + if (use.getOwner() == consumer) + use.set(supData); + } + } + locGraph->containerOp->erase(); +} diff --git a/experimental/lib/Transforms/HandshakeStraightToQueue.cpp b/experimental/lib/Transforms/HandshakeStraightToQueue.cpp index c276561d4b..e68ef707c4 100644 --- a/experimental/lib/Transforms/HandshakeStraightToQueue.cpp +++ b/experimental/lib/Transforms/HandshakeStraightToQueue.cpp @@ -17,16 +17,21 @@ #include "dynamatic/Support/DynamaticPass.h" #include "experimental/Support/CFGAnnotation.h" #include "experimental/Support/FtdImplementation.h" +#include "mlir/Dialect/Arith/IR/Arith.h" #include "mlir/Dialect/ControlFlow/IR/ControlFlow.h" +#include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h" #include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Support/LogicalResult.h" #include "mlir/Transforms/DialectConversion.h" +#include "llvm/Support/Debug.h" #include using namespace mlir; using namespace dynamatic; using namespace dynamatic::experimental; +#define DEBUG_TYPE "handshake-straight-to-queue" + // [START Boilerplate code for the MLIR pass] #include "experimental/Transforms/Passes.h" // IWYU pragma: keep namespace dynamatic { @@ -300,33 +305,66 @@ static void minimizeGroupsConnections(handshake::FuncOp funcOp, /// For each element in the group, build a lazy fork and use its output to feed /// the correspondent input of the LSQ. -static DenseMap +static handshake::LazyForkOp +ensureLazyForkOutputCount(Block *bb, Value input, unsigned numResults, + DenseMap &forksGraph, + PatternRewriter &rewriter) { + auto it = forksGraph.find(bb); + if (it == forksGraph.end()) { + rewriter.setInsertionPointToStart(bb); + auto forkOp = rewriter.create(bb->front().getLoc(), + input, numResults); + forksGraph[bb] = forkOp; + return forkOp; + } + + handshake::LazyForkOp oldFork = it->second; + if (oldFork->getNumResults() >= numResults) + return oldFork; + + rewriter.setInsertionPoint(oldFork); + auto newFork = rewriter.create( + oldFork.getLoc(), oldFork.getOperand(), numResults); + newFork->setAttrs(oldFork->getAttrs()); + + for (unsigned idx = 0, e = oldFork->getNumResults(); idx < e; ++idx) + oldFork->getResult(idx).replaceAllUsesWith(newFork->getResult(idx)); + + rewriter.eraseOp(oldFork); + forksGraph[bb] = newFork; + return newFork; +} + +static void connectLSQToForkGraph(handshake::FuncOp &funcOp, DenseSet &groups, handshake::LSQOp lsqOp, + DenseMap &sharedForks, + DenseMap &nextAllocPort, PatternRewriter &rewriter) { - DenseMap forksGraph; + DenseMap allocPort; auto startValue = (Value)funcOp.getArguments().back(); - // Create the fork nodes: for each group among the set of groups + // Create or reuse one fork node per basic block. Result #0 is reserved for + // the continuation/dependency network; one extra result is allocated for each + // LSQ group allocation control input in the block. for (MemoryGroup *group : groups) { Block *bb = group->bb; - rewriter.setInsertionPointToStart(bb); - // Add a lazy fork with two outputs, having the start control value as - // input and two output ports, one for the LSQ and one for the subsequent - // buffer - auto forkOp = rewriter.create(bb->front().getLoc(), - startValue, 2); + unsigned port = nextAllocPort.lookup(bb); + if (port == 0) + port = 1; + nextAllocPort[bb] = port + 1; + allocPort[bb] = port; - // Add the new component to the list of components create for FTD and to - // the fork graph - forksGraph[bb] = forkOp; + auto forkOp = ensureLazyForkOutputCount(bb, startValue, port + 1, + sharedForks, rewriter); + sharedForks[bb] = forkOp; } - // The second output of each lazy fork must be connected to the LSQ, so that - // they can activate the allocation for the operations of the corresponding - // basic block + // The allocated output of each lazy fork must be connected to the LSQ, so + // that it can activate the allocation for the operations of the corresponding + // basic block. // // For each input of the LSQ for (auto [opIdx, op] : llvm::enumerate(lsqOp.getOperands())) { @@ -338,12 +376,10 @@ connectLSQToForkGraph(handshake::FuncOp &funcOp, // fork of the graph auto cmerge = llvm::dyn_cast(op.getDefiningOp()); Block *bb = cmerge->getBlock(); - if (!forksGraph.contains(bb)) + if (!allocPort.contains(bb)) continue; - lsqOp.setOperand(opIdx, forksGraph[bb]->getResult(1)); + lsqOp.setOperand(opIdx, sharedForks[bb]->getResult(allocPort[bb])); } - - return forksGraph; } /// Get all the load and store operations related to a LSQ operation @@ -363,27 +399,28 @@ getLsqOps(handshake::FuncOp &funcOp, handshake::LSQOp lsqOp) { /// Given a graph of lazy forks, connect the elements together with some proper /// SSA phi static LogicalResult -connectForkGraph(handshake::FuncOp &funcOp, - const DenseSet &groupsGraph, - const DenseMap &forksGraph, - PatternRewriter &rewriter) { - - for (MemoryGroup *consumerGroup : groupsGraph) { - - DenseMap> deps; - SmallVector forkDeps; - - for (auto &producerGroup : consumerGroup->preds) { - Operation *producerLF = forksGraph.at(producerGroup->bb); - forkDeps.push_back(producerLF->getResult(0)); +connectForkGraphs(handshake::FuncOp &funcOp, + ArrayRef> allGroupsGraphs, + const DenseMap &forksGraph, + PatternRewriter &rewriter) { + + DenseMap> deps; + + for (const DenseSet &groupsGraph : allGroupsGraphs) { + for (MemoryGroup *consumerGroup : groupsGraph) { + handshake::LazyForkOp consumerLF = forksGraph.at(consumerGroup->bb); + SmallVector &forkDeps = deps[&consumerLF->getOpOperand(0)]; + + for (auto &producerGroup : consumerGroup->preds) { + Operation *producerLF = forksGraph.at(producerGroup->bb); + Value producerToken = producerLF->getResult(0); + if (!llvm::is_contained(forkDeps, producerToken)) + forkDeps.push_back(producerToken); + } } - - deps[&forksGraph.at(consumerGroup->bb)->getOpOperand(0)] = forkDeps; - - if (failed(ftd::createPhiNetworkDeps(funcOp.getRegion(), rewriter, deps))) - return failure(); } - return success(); + + return ftd::createPhiNetworkDeps(funcOp.getRegion(), rewriter, deps); } /// Remove the network of cmerges in case the function is void. The SQ pass @@ -475,6 +512,176 @@ static void removeNetworkCMerges(handshake::FuncOp &funcOp, toRemove->erase(); } +/// Per-block edge information captured from the multi-block CFG before +/// flattening. Used to reconstruct a ShadowCFG afterwards. +struct CapturedEdgeInfo { + bool isConditional = false; + bool hasSuccessors = false; + unsigned trueSuccIdx = 0; + unsigned falseSuccIdx = 0; + unsigned uncondSuccIdx = 0; +}; + +/// Capture the CFG topology and branch condition Values of a handshake::FuncOp +/// while it still has restored multi-block structure. Both pieces of +/// information are destroyed by removeNetworkCMerges + flattenFunction and are +/// needed later to construct the ShadowCFG for addRegen / addSupp. +static void captureCFGTopology(handshake::FuncOp funcOp, unsigned &numBlocks, + SmallVector &edges, + DenseMap &capturedConditions) { + Region ®ion = funcOp.getRegion(); + DenseMap blockToIdx; + for (auto [idx, block] : llvm::enumerate(region)) + blockToIdx[&block] = idx; + + numBlocks = blockToIdx.size(); + edges.resize(numBlocks); + + for (auto [idx, block] : llvm::enumerate(region)) { + CapturedEdgeInfo &edge = edges[idx]; + Operation *term = block.getTerminator(); + if (auto condBr = dyn_cast(term)) { + edge.isConditional = true; + edge.hasSuccessors = true; + edge.trueSuccIdx = blockToIdx.lookup(condBr.getTrueDest()); + edge.falseSuccIdx = blockToIdx.lookup(condBr.getFalseDest()); + } else if (auto br = dyn_cast(term)) { + edge.hasSuccessors = true; + edge.uncondSuccIdx = blockToIdx.lookup(br.getDest()); + } + } + + // Capture branch condition Values from handshake::ConditionalBranchOp ops. + // This must happen before removeNetworkCMerges, which erases cmerge-network + // ConditionalBranchOps and may leave some blocks without any surviving + // condition source. We prefer conditions that do NOT originate from a + // ControlMergeOp (those belong to the cmerge network and will be erased). + funcOp.walk([&](handshake::ConditionalBranchOp brOp) { + auto bbAttr = brOp->getAttrOfType("handshake.bb"); + if (!bbAttr) + return; + unsigned bbIdx = bbAttr.getUInt(); + Value cond = brOp.getConditionOperand(); + bool fromCmerge = + cond.getDefiningOp() && + llvm::isa(cond.getDefiningOp()); + // Always prefer a non-cmerge condition; keep the first non-cmerge one found + if (!capturedConditions.contains(bbIdx)) { + // First condition seen for this BB — always capture + capturedConditions[bbIdx] = cond; + } else if (!fromCmerge) { + // Found a better (non-cmerge) condition — upgrade + capturedConditions[bbIdx] = cond; + } + }); +} + +/// Build a ShadowCFG for the given handshake::FuncOp using previously captured +/// CFG topology and condition Values. The shadow function replicates the +/// original block structure with standard CF terminators, enabling analysis +/// passes (BlockIndexing, DominanceInfo, CFGLoopInfo) that addRegen and addSupp +/// rely on. +static ftd::ShadowCFG buildShadowFromCapturedTopology( + OpBuilder &builder, handshake::FuncOp funcOp, unsigned numBlocks, + const SmallVector &edges, + const DenseMap &capturedConditions) { + + ftd::ShadowCFG shadow; + Location loc = funcOp.getLoc(); + + // Create a temporary func::FuncOp with blocks + CF terminators that mirror + // the original multi-block topology. + { + OpBuilder::InsertionGuard guard(builder); + builder.setInsertionPointAfter(funcOp); + auto funcType = builder.getFunctionType({}, {}); + shadow.shadowFunc = + builder.create(loc, "__ftd_shadow_cfg__", funcType); + + Region &R = shadow.shadowFunc.getBody(); + SmallVector blocks; + for (unsigned i = 0; i < numBlocks; ++i) + blocks.push_back(builder.createBlock(&R, R.end())); + + for (unsigned i = 0; i < numBlocks; ++i) { + const CapturedEdgeInfo &edge = edges[i]; + builder.setInsertionPointToEnd(blocks[i]); + + if (edge.isConditional) { + auto dummyCond = + builder.create(loc, builder.getBoolAttr(true)); + builder.create( + loc, dummyCond, blocks[edge.trueSuccIdx], ValueRange{}, + blocks[edge.falseSuccIdx], ValueRange{}); + } else if (edge.hasSuccessors) { + builder.create(loc, blocks[edge.uncondSuccIdx]); + } else { + builder.create(loc); + } + } + } + + // Use the pre-captured condition Values (captured before removeNetworkCMerges + // destroyed cmerge-network ConditionalBranchOps). + shadow.conditionMap = capturedConditions; + + return shadow; +} + +/// Remove the cmerge network, flatten the function, and run FTD gating on the +/// flattened IR using a temporary ShadowCFG reconstructed from the original +/// multi-block topology. +static LogicalResult runPostCmergeFtd(handshake::FuncOp funcOp, + MLIRContext *ctx, + bool resolveCondPlaceholders) { + ConversionPatternRewriter rewriter(ctx); + + experimental::cfg::markBasicBlocks(funcOp, rewriter); + + // Capture the original CFG before the cmerge network is removed so we can + // rebuild a ShadowCFG after flattening for addRegen/addSupp. + unsigned capturedNumBlocks = 0; + SmallVector capturedEdges; + DenseMap capturedConditions; + captureCFGTopology(funcOp, capturedNumBlocks, capturedEdges, + capturedConditions); + + removeNetworkCMerges(funcOp, rewriter); + + if (failed(cfg::flattenFunction(funcOp))) + return failure(); + + if (capturedNumBlocks <= 1) + return success(); + + OpBuilder builder(ctx); + ftd::ShadowCFG shadow = buildShadowFromCapturedTopology( + builder, funcOp, capturedNumBlocks, capturedEdges, capturedConditions); + + if (resolveCondPlaceholders) { + ftd::resolveCondPlaceholders(funcOp, builder, shadow); + + // Pick up negated condition signals produced by resolveCondPlaceholders. + for (auto notOp : funcOp.getOps()) { + if (!notOp->hasAttr("ftd.cvar")) + continue; + auto bbAttr = notOp->getAttrOfType("handshake.bb"); + if (!bbAttr) + continue; + shadow.conditionMap[bbAttr.getUInt()] = notOp.getResult(); + } + } + + ftd::addRegen(funcOp, builder, shadow); + ftd::addSupp(funcOp, builder, shadow); + + if (resolveCondPlaceholders) + ftd::finalizeCondPlaceholders(funcOp); + + shadow.destroy(); + return success(); +} + /// Run straight to the queue. static LogicalResult applyStraightToQueue(handshake::FuncOp funcOp, MLIRContext *ctx) { @@ -483,14 +690,19 @@ static LogicalResult applyStraightToQueue(handshake::FuncOp funcOp, // Return if there are no LSQs in the function if (funcOp.getOps().empty()) { - removeNetworkCMerges(funcOp, rewriter); - return success(); + if (failed(cfg::restoreCfStructure(funcOp, rewriter))) + return failure(); + return runPostCmergeFtd(funcOp, ctx, /*resolveCondPlaceholders=*/false); } // Restore the cf structure to work on a structured IR if (failed(cfg::restoreCfStructure(funcOp, rewriter))) return failure(); + SmallVector> allGroupsGraphs; + DenseMap sharedForks; + DenseMap nextAllocPort; + // For each LSQ for (const handshake::LSQOp lsqOp : funcOp.getOps()) { @@ -500,8 +712,7 @@ static LogicalResult applyStraightToQueue(handshake::FuncOp funcOp, // Get all the memory depdencies among the operations connected to the // same LSQ auto lsqMemDeps = identifyMemoryDependencies(funcOp, lsqOps); - for (auto &dep : lsqMemDeps) - dep.print(); + LLVM_DEBUG(for (auto &dep : lsqMemDeps) dep.print()); // Build a group graph out of the dependencies auto groupsGraph = constructGroupsGraph(lsqOps, lsqMemDeps); @@ -509,41 +720,37 @@ static LogicalResult applyStraightToQueue(handshake::FuncOp funcOp, // Apply group minimization techniques minimizeGroupsConnections(funcOp, groupsGraph); - for (auto &g : groupsGraph) - g->print(); + LLVM_DEBUG(for (auto &g : groupsGraph) g->print()); // Build a lazy fork for each group and connect it to the related // activation input in the LSQ - auto forksGraph = - connectLSQToForkGraph(funcOp, groupsGraph, lsqOp, rewriter); - - // Connect the lazy forks together through a network of merges - if (failed(connectForkGraph(funcOp, groupsGraph, forksGraph, rewriter))) - return failure(); + connectLSQToForkGraph(funcOp, groupsGraph, lsqOp, sharedForks, + nextAllocPort, rewriter); - // Delete the groups - for (auto *g : groupsGraph) - delete g; + allGroupsGraphs.push_back(std::move(groupsGraph)); } - // Replace each merge created by `createPhiNetwork` with a multiplxer - if (failed(ftd::replaceMergeToGSA(funcOp, rewriter))) + // Connect the shared lazy forks together through a network of merges. + if (failed(connectForkGraphs(funcOp, allGroupsGraphs, sharedForks, rewriter))) return failure(); - // Run fast token delivery on the newly inserted operations - experimental::ftd::addRegen(funcOp, rewriter); - experimental::ftd::addSupp(funcOp, rewriter); - experimental::cfg::markBasicBlocks(funcOp, rewriter); + // Delete the groups + for (const DenseSet &groupsGraph : allGroupsGraphs) + for (auto *g : groupsGraph) + delete g; - // Try to remove the network of cmerges if possible (i.e. if the function was - // void) - removeNetworkCMerges(funcOp, rewriter); + // Create condition placeholders for every conditional block in the restored + // CFG. The refactored addGsaGates (called by replaceMergeToGSA) relies on + // these placeholders when building distribution networks and BDD circuits + // for MU gates. In the FTD flow, createAllCondPlaceholders is called before + // addGsaGates; S2Q must do the same. + ftd::createAllCondPlaceholders(funcOp.getRegion(), rewriter); - // Remove the blocks and terminators - if (failed(cfg::flattenFunction(funcOp))) + // Replace each merge created by `createPhiNetwork` with a multiplxer + if (failed(ftd::replaceMergeToGSA(funcOp, rewriter))) return failure(); - return success(); + return runPostCmergeFtd(funcOp, ctx, /*resolveCondPlaceholders=*/true); } struct HandshakeStraightToQueuePass diff --git a/include/dynamatic/Dialect/Handshake/HandshakeOps.td b/include/dynamatic/Dialect/Handshake/HandshakeOps.td index d39529a5d3..f957bf1fba 100644 --- a/include/dynamatic/Dialect/Handshake/HandshakeOps.td +++ b/include/dynamatic/Dialect/Handshake/HandshakeOps.td @@ -341,10 +341,6 @@ def InitOp : Handshake_Op<"init", [ let arguments = (ins HandshakeType:$operand); let results = (outs HandshakeType:$result); - let extraClassDeclaration = [{ - static constexpr ::llvm::StringLiteral INIT_TOKEN_ATTR_NAME = "INIT_TOKEN", - TIMING_ATTR_NAME = "TIMING"; - }]; let assemblyFormat = [{ $operand attr-dict `:` custom(type($operand)) diff --git a/include/dynamatic/Support/CFG.h b/include/dynamatic/Support/CFG.h index 0ef76046c2..c275f31860 100644 --- a/include/dynamatic/Support/CFG.h +++ b/include/dynamatic/Support/CFG.h @@ -98,12 +98,12 @@ bool getBBEndpoints(Value val, Operation *user, BBEndpoints &endpoints); bool getBBEndpoints(Value val, BBEndpoints &endpoints); /// Determines whether the value is a backedge i.e., whether the channel -/// corresponding to the value is located between a branch-like operation and a -/// merge-like operation, where the merge-like operation happens semantically -/// "before" the branch-like operation. This function can only correctly -/// identify backedges if the circuit's branches and merges are associated to -/// basic blocks (otherwise it will always return false). `user` must be one of -/// `val`'s users. +/// corresponding to the value is located between a loop-feedback source +/// (branch-like or compare-like within a block) and a merge-like operation, +/// where the merge-like operation happens semantically "before" that source. +/// This function can only correctly identify backedges if the circuit's +/// branches, compares, and merges are associated to basic blocks (otherwise it +/// will always return false). `user` must be one of `val`'s users. bool isBackedge(Value val, Operation *user, BBEndpoints *endpoints = nullptr); /// Determines whether the value is a backedge. The value must have a single diff --git a/include/dynamatic/Transforms/BufferPlacement/Utils/CFDFC.h b/include/dynamatic/Transforms/BufferPlacement/Utils/CFDFC.h index 186c3aaab3..1df3872521 100644 --- a/include/dynamatic/Transforms/BufferPlacement/Utils/CFDFC.h +++ b/include/dynamatic/Transforms/BufferPlacement/Utils/CFDFC.h @@ -24,6 +24,7 @@ #include "dynamatic/Support/ConstraintProgramming/ConstraintProgramming.h" #include "dynamatic/Support/LLVM.h" #include "experimental/Support/StdProfiler.h" +#include "mlir/Support/LogicalResult.h" namespace dynamatic { namespace buffer { @@ -59,13 +60,20 @@ struct CFDFC { /// Constructs a CFDFC from a set of selected archs and basic blocks in the /// function. Assumes that every value in the function is used exactly once. - CFDFC(handshake::FuncOp funcOp, ArchSet &archs, unsigned numExec); + CFDFC(handshake::FuncOp funcOp, ArchSet &archs, unsigned numExec, + DenseSet backwardChannels); // Determines whether the channel is a "CFDFC backedge" i.e., the first // channel along a sequence of backedges from a source block to a destination // block. The distinction is important for the buffer placement MILP, which // uses backedges to determine where to insert "tokens" in the circuit. static bool isCFDFCBackedge(Value val); + + // Determines whether the channel has a corresponding CFG edge. + bool isCFGCompliant(unsigned srcBB, unsigned dstBB); + + // AYA: Added this from Jiahui to write the CFDFCs in DOT for debugging + void writeDot(const std::string &fileName); }; /// Represents a union of CFDFCs. Its blocks, units, channels, and backedges are diff --git a/lib/Analysis/ControlDependenceAnalysis.cpp b/lib/Analysis/ControlDependenceAnalysis.cpp index 492ab1345e..2c4fb66142 100644 --- a/lib/Analysis/ControlDependenceAnalysis.cpp +++ b/lib/Analysis/ControlDependenceAnalysis.cpp @@ -124,10 +124,6 @@ void dynamatic::ControlDependenceAnalysis::identifyAllControlDeps( Block *leastCommonAnc = postDomInfo.findNearestCommonDominator(successor, &bb); - // Loop case - if (leastCommonAnc == &bb) - blocksControlDeps[&bb].allControlDeps.insert(&bb); - // In the post dominator tree, all the nodes from `leastCommonAnc` to // `successor` should be control dependent on `block` blocksControlDeps[successor].allControlDeps.insert(&bb); @@ -159,14 +155,24 @@ void dynamatic::ControlDependenceAnalysis::identifyAllControlDeps( void dynamatic::ControlDependenceAnalysis::addDepsOfDeps(Region ®ion) { - // For each block, consider each of its dependencies (`oneDep`) and move each - // of its dependencies into block's - for (Block &block : region.getBlocks()) { - BlockControlDeps blockControlDeps = blocksControlDeps[&block]; - for (auto &oneDep : blockControlDeps.allControlDeps) { - DenseSet &oneDepDeps = blocksControlDeps[oneDep].allControlDeps; - for (auto &oneDepDep : oneDepDeps) - blocksControlDeps[&block].allControlDeps.insert(oneDepDep); + bool changed = true; + while (changed) { + changed = false; + + // For each block, consider each of its dependencies and move each + // of its dependencies into block's + for (Block &block : region.getBlocks()) { + DenseSet &blockDeps = blocksControlDeps[&block].allControlDeps; + SmallVector currentDeps(blockDeps.begin(), blockDeps.end()); + + for (Block *oneDep : currentDeps) { + DenseSet &oneDepDeps = + blocksControlDeps[oneDep].allControlDeps; + for (Block *oneDepDep : oneDepDeps) { + if (blockDeps.insert(oneDepDep).second) + changed = true; // Found a new dependency + } + } } } } diff --git a/lib/Conversion/CfToHandshake/CfToHandshake.cpp b/lib/Conversion/CfToHandshake/CfToHandshake.cpp index 54a6331a04..7d55ce71aa 100644 --- a/lib/Conversion/CfToHandshake/CfToHandshake.cpp +++ b/lib/Conversion/CfToHandshake/CfToHandshake.cpp @@ -1194,6 +1194,8 @@ LogicalResult ConvertIndexCast::matchAndRewrite( ConversionPatternRewriter &rewriter) const { auto getWidth = [](Type type) -> unsigned { + if (auto chanTy = dyn_cast(type)) + type = chanTy.getDataType(); if (isa(type)) return 32; return type.getIntOrFloatBitWidth(); @@ -1504,13 +1506,13 @@ ConvertConstants::matchAndRewrite(arith::ConstantOp cstOp, // Determine the new constant's control input Value controlVal; - if (isCstSourcable(cstOp)) { - auto sourceOp = rewriter.create(cstOp.getLoc()); - inheritBB(cstOp, sourceOp); - controlVal = sourceOp.getResult(); - } else { - controlVal = getBlockControl(cstOp); - } + // if (isCstSourcable(cstOp)) { + // auto sourceOp = rewriter.create(cstOp.getLoc()); + // inheritBB(cstOp, sourceOp); + // controlVal = sourceOp.getResult(); + // } else { + controlVal = getBlockControl(cstOp); + //} TypedAttr cstAttr = cstOp.getValue(); // Convert IndexType'd values to equivalent signless integers diff --git a/lib/Conversion/HandshakeToHW/HandshakeToHW.cpp b/lib/Conversion/HandshakeToHW/HandshakeToHW.cpp index 10ea0d80ad..fa0dbde914 100644 --- a/lib/Conversion/HandshakeToHW/HandshakeToHW.cpp +++ b/lib/Conversion/HandshakeToHW/HandshakeToHW.cpp @@ -768,6 +768,18 @@ ModuleDiscriminator::ModuleDiscriminator(Operation *op) { addUnsigned("DATA_WIDTH", resType.getElementTypeBitWidth()); addUnsigned("SIZE", resType.getNumElements()); }) + .Case([&](handshake::InitOp initOp) { + auto paramsAttr = + initOp->getAttrOfType("hw.parameters"); + if (paramsAttr) { + auto initTokenAttr = + paramsAttr.get("INIT_TOKEN").dyn_cast_or_null(); + int initialValue = + (initTokenAttr && initTokenAttr.getValue()) ? 1 : 0; + addUnsigned("INITIAL_VALUE", initialValue); + } else + addUnsigned("INITIAL_VALUE", 0); + }) .Default([&](auto) { op->emitError() << "This operation cannot be lowered to RTL " "due to a lack of an RTL implementation for it."; @@ -2181,6 +2193,7 @@ class HandshakeToHWPass ConvertToHWInstance, ConvertToHWInstance, ConvertToHWInstance, + ConvertToHWInstance, ConvertToHWInstance, ConvertToHWInstance, ConvertToHWInstance, diff --git a/lib/Support/CFG.cpp b/lib/Support/CFG.cpp index e05541ef4d..fd45884213 100644 --- a/lib/Support/CFG.cpp +++ b/lib/Support/CFG.cpp @@ -150,23 +150,31 @@ static bool followToBlock(Operation *op, unsigned &bb, } /// Determines whether the operation is of a nature which can be traversed -/// outside blocks during backedge identification. -static inline bool canGoThroughOutsideBlocks(Operation *op) { +/// backwards during backedge source identification. +static inline bool canBacktrackToLoopSourceThrough(Operation *op) { return isa(op); } -/// Attempts to backtrack through forks and bitwidth modification operations -/// till reaching a branch-like operation. On success, returns the branch-like -/// operation that was backtracked to (or the passed operation if it was itself -/// branch-like); otherwise, returns nullptr. -static Operation *backtrackToBranch(Operation *op) { +/// Determines whether the operation is of a nature which can be traversed +/// forwards during backedge destination identification. +static inline bool canFollowToMergeThrough(Operation *op) { + return isa(op); +} + +/// Attempts to backtrack through source-transparent operations till reaching +/// an operation that can act as the source of loop feedback within a block. On +/// success, returns that operation (or the passed operation if it was itself +/// such a source); otherwise, returns nullptr. +static Operation *backtrackToLoopSource(Operation *op) { do { if (!op) break; - if (isa(op)) + if (isa(op)) return op; - if (canGoThroughOutsideBlocks(op)) + if (canBacktrackToLoopSourceThrough(op)) op = op->getOperand(0).getDefiningOp(); else break; @@ -175,14 +183,14 @@ static Operation *backtrackToBranch(Operation *op) { } /// Attempts to follow the def-use chains of all the operation's results through -/// forks and bitwidth modification operations till reaching merge-like -/// operations that all belong to the same basic block. On success, returns one -/// of the merge-like operations reached by a def-use chain (or the passed -/// operation if it was itself merge-like); otherwise, returns nullptr. +/// destination-transparent operations till reaching merge-like operations that +/// all belong to the same basic block. On success, returns one of the +/// merge-like operations reached by a def-use chain (or the passed operation if +/// it was itself merge-like); otherwise, returns nullptr. static Operation *followToMerge(Operation *op) { if (isa(op)) return op; - if (canGoThroughOutsideBlocks(op)) { + if (canFollowToMergeThrough(op)) { // All users of the operation's results must lead to merges within a unique // block SmallVector mergeOps; @@ -245,20 +253,20 @@ bool dynamatic::isBackedge(Value val, Operation *user, BBEndpoints *endpoints) { return false; // If both source and destination blocks are identical, the edge must be - // located between a branch-like operation and a merge-like operation - Operation *brOp = backtrackToBranch(val.getDefiningOp()); - if (!brOp) + // located between a loop-feedback source and a merge-like operation. + Operation *srcOp = backtrackToLoopSource(val.getDefiningOp()); + if (!srcOp) return false; Operation *mergeOp = followToMerge(user); if (!mergeOp) return false; - // Check that the branch and merge are part of the same block indicated by the + // Check that the source and merge are part of the same block indicated by the // edge's BB endpoints (should be the case in all non-degenerate cases) - std::optional brBB = getLogicBB(brOp); + std::optional srcBB = getLogicBB(srcOp); std::optional mergeBB = getLogicBB(mergeOp); - return brBB.has_value() && mergeBB.has_value() && *brBB == *mergeBB && - *brBB == bbs.srcBB; + return srcBB.has_value() && mergeBB.has_value() && *srcBB == *mergeBB && + *srcBB == bbs.srcBB; } bool dynamatic::isBackedge(Value val, BBEndpoints *endpoints) { diff --git a/lib/Support/RTL/RTL.cpp b/lib/Support/RTL/RTL.cpp index 47a4ae33bf..7b8d80d219 100644 --- a/lib/Support/RTL/RTL.cpp +++ b/lib/Support/RTL/RTL.cpp @@ -337,7 +337,8 @@ LogicalResult RTLMatch::registerBitwidthParameter(hw::HWModuleExternOp &modOp, handshakeOp == "handshake.spec_commit" || handshakeOp == "handshake.spec_save_commit" || handshakeOp == "handshake.sharing_wrapper" || - handshakeOp == "handshake.non_spec" + handshakeOp == "handshake.non_spec" || + handshakeOp == "handshake.init" // clang-format on ) { // Default @@ -493,7 +494,8 @@ RTLMatch::registerExtraSignalParameters(hw::HWModuleExternOp &modOp, handshakeOp == "handshake.load" || handshakeOp == "handshake.store" || handshakeOp == "handshake.spec_commit" || - handshakeOp == "handshake.speculating_branch" + handshakeOp == "handshake.speculating_branch" || + handshakeOp == "handshake.init" // clang-format on ) { diff --git a/lib/Transforms/BufferPlacement/CMakeLists.txt b/lib/Transforms/BufferPlacement/CMakeLists.txt index 8af26da499..8f1734c1a9 100644 --- a/lib/Transforms/BufferPlacement/CMakeLists.txt +++ b/lib/Transforms/BufferPlacement/CMakeLists.txt @@ -36,3 +36,13 @@ if (GUROBI_FOUND) ) endif() +# Find the libraries from Graphviz +find_package(PkgConfig REQUIRED) +pkg_check_modules(GVC REQUIRED libgvc) +pkg_check_modules(CGRAPH REQUIRED libcgraph) + +TARGET_LINK_LIBRARIES(DynamaticBufferPlacement + LINK_PUBLIC + ${GVC_LIBRARIES} ${CGRAPH_LIBRARIES} +) + diff --git a/lib/Transforms/BufferPlacement/HandshakePlaceBuffers.cpp b/lib/Transforms/BufferPlacement/HandshakePlaceBuffers.cpp index e8230b7cf0..de6434fd2f 100644 --- a/lib/Transforms/BufferPlacement/HandshakePlaceBuffers.cpp +++ b/lib/Transforms/BufferPlacement/HandshakePlaceBuffers.cpp @@ -36,8 +36,14 @@ #include "llvm/Support/Casting.h" #include "llvm/Support/Path.h" #include +#include #include +#include "mlir/Support/LLVM.h" +#include "llvm/ADT/DenseMap.h" +#include "llvm/ADT/DenseSet.h" +#include "llvm/ADT/StringSet.h" + using namespace mlir; using namespace dynamatic; using namespace dynamatic::handshake; @@ -208,6 +214,8 @@ LogicalResult HandshakePlaceBuffersPass::placeUsingMILP() { } ModuleOp modOp = llvm::dyn_cast(getOperation()); + + // AYA: TODO: This is where I should add the timing of the OOE units // // Read the operations' timing models from disk TimingDatabase timingDB; @@ -228,9 +236,10 @@ LogicalResult HandshakePlaceBuffersPass::placeUsingMILP() { << "Failed to read profiling information from CSV"; } + // AYA commented this out // Check IR invariants and parse basic block archs from disk - if (failed(checkFuncInvariants(info))) - return failure(); + // if (failed(checkFuncInvariants(info))) + // return failure(); // Get CFDFCs from the function unless the functions has no archs (i.e., // it has a single block) in which case there are no CFDFCs @@ -377,11 +386,208 @@ static void logFuncInfo(FuncInfo &info) { os << "- Number of channels: " << cf->channels.size() << "\n"; os << "- Number of backedges: " << cf->backedges.size() << "\n\n"; os.unindent(); + + // AYA: Added this from Jiahui to write the CFDFCs in DOT for debugging + cf->writeDot("AYAA-cfdfc_" + std::to_string(idx) + ".dot"); } os.flush(); } +////////////////////////////////////////////////////////////////////////////////////////// +// AYA: The following functions are for identifying backward edges in the +// circuit graph +static void +printBackwardChannels(const llvm::DenseSet &backwardChannels) { + llvm::errs() << "=== Backward Channels (with src/dst ops) ===\n"; + int idx = 0; + + for (Value v : backwardChannels) { + llvm::errs() << "Channel " << idx++ << ":\n"; + + // Source operation + if (Operation *srcOp = v.getDefiningOp()) { + llvm::errs() << " Source: "; + srcOp->print(llvm::errs()); + llvm::errs() << "\n"; + } else { + llvm::errs() << " Source: \n"; + } + + // Destination operations + for (auto &use : v.getUses()) { + Operation *dstOp = use.getOwner(); + llvm::errs() << " Destination: "; + dstOp->print(llvm::errs()); + llvm::errs() << "\n"; + } + llvm::errs() << "\n"; + } + + if (backwardChannels.empty()) + llvm::errs() << "(empty)\n"; +} + +// Aya: The following captivates the necessary logic for extracting back edges +// in cycles +namespace { + +struct CircuitEdge { + Operation *src; + Operation *dst; + Value channel; +}; + +static bool isBackedgeSourceLike(Operation *op) { + do { + if (!op) + return false; + if (isa(op)) + return true; + if (isa(op)) + op = op->getOperand(0).getDefiningOp(); + else + return false; + } while (true); +} + +static bool isBackedgeDestinationLike(Operation *op) { + if (isa(op)) + return true; + + auto notOp = dyn_cast(op); + if (!notOp) + return false; + + return llvm::any_of(notOp.getResult().getUsers(), [](Operation *user) { + return isa(user); + }); +} + +/// Finds all loop-feedback-source -> merge-like backward channels per cyclic +/// SCC in the handshake graph. Grouping by SCC remains more stable than trying +/// to assign channels to every simple cycle when cycles overlap. +static mlir::DenseSet +findBackwardChannelPerCyclicRegion(handshake::FuncOp funcOp) { + SmallVector ops; + SmallVector edges; + llvm::DenseMap> succs; + + for (Operation &op : funcOp.getOps()) { + ops.push_back(&op); + succs[&op] = {}; + } + + for (Operation *src : ops) { + for (Value result : src->getResults()) { + for (OpOperand &use : result.getUses()) { + Operation *dst = use.getOwner(); + + if (isa(dst)) + continue; + + edges.push_back({src, dst, result}); + succs[src].push_back(dst); + } + } + } + + llvm::DenseMap index, lowlink; + llvm::DenseSet onStack; + SmallVector stack; + SmallVector> sccs; + unsigned nextIndex = 0; + + std::function strongConnect = [&](Operation *op) { + index[op] = nextIndex; + lowlink[op] = nextIndex; + ++nextIndex; + stack.push_back(op); + onStack.insert(op); + + for (Operation *succ : succs[op]) { + if (!index.contains(succ)) { + strongConnect(succ); + lowlink[op] = std::min(lowlink[op], lowlink[succ]); + } else if (onStack.contains(succ)) { + lowlink[op] = std::min(lowlink[op], index[succ]); + } + } + + if (lowlink[op] != index[op]) + return; + + SmallVector scc; + while (true) { + Operation *top = stack.pop_back_val(); + onStack.erase(top); + scc.push_back(top); + if (top == op) + break; + } + sccs.push_back(std::move(scc)); + }; + + for (Operation *op : ops) { + if (!index.contains(op)) + strongConnect(op); + } + + mlir::DenseSet backwardChannels; + for (const auto &scc : sccs) { + llvm::DenseSet sccNodes(scc.begin(), scc.end()); + + bool isCyclic = scc.size() > 1; + if (!isCyclic) { + Operation *only = scc.front(); + isCyclic = llvm::any_of(edges, [&](const CircuitEdge &edge) { + return edge.src == only && edge.dst == only; + }); + } + if (!isCyclic) + continue; + + for (const CircuitEdge &edge : edges) { + if (!sccNodes.contains(edge.src) || !sccNodes.contains(edge.dst)) + continue; + if (!isBackedge(edge.channel)) + continue; + if (!isBackedgeSourceLike(edge.src)) + continue; + if (!isBackedgeDestinationLike(edge.dst)) + continue; + backwardChannels.insert(edge.channel); + } + } + + return backwardChannels; +} + +} // namespace + +// useful in debugging +// static void printCycles(const CycleList &cycles) { +// llvm::errs() << "=== Circuit Cycles ===\n"; +// int cycleIdx = 0; + +// for (const auto &cycle : cycles) { +// llvm::errs() << "Cycle " << cycleIdx++ << " (" << cycle.size() +// << " ops):\n"; + +// for (Operation *op : cycle) { +// llvm::errs() << " - "; +// op->print(llvm::errs()); +// llvm::errs() << "\n"; +// } + +// llvm::errs() << "\n"; +// } +// } + +//////////////////////////////////////////////////////////////////////////////////// + LogicalResult HandshakePlaceBuffersPass::getCFDFCs(FuncInfo &info, std::vector &cfdfcs) { SmallVector archsCopy(info.archs); @@ -400,6 +606,25 @@ LogicalResult HandshakePlaceBuffersPass::getCFDFCs(FuncInfo &info, bbs.insert(arch.dstBB); } + //////////// AYA added the following to identify all graph cycles in the + /// circuit graph + // Identify the cycles and the backward channels in your circuit + llvm::errs() << "\nBefore findALlCycles\n"; + + // CycleList circuitCycles = + // findAllCycles(info.funcOp); + // // llvm::errs() << "\nAfter findALlCycles\n"; + + // AYA TO AYA: The goal is to send mlir::DenseSet backwardChannels + // structure to the constructor of CFDFC below. + // GraphForJohnson johnsonGraph(info.funcOp); + // CycleList circuitCycles = johnsonGraph.findAllCycles(); + + // printCycles(circuitCycles); + mlir::DenseSet backwardChannels = + findBackwardChannelPerCyclicRegion(info.funcOp); + printBackwardChannels(backwardChannels); + // Set of selected archs ArchSet selectedArchs; // Number of executions @@ -424,7 +649,7 @@ LogicalResult HandshakePlaceBuffersPass::getCFDFCs(FuncInfo &info, break; // Create the CFDFC from the set of selected archs and BBs - cfdfcs.emplace_back(info.funcOp, selectedArchs, numExecs); + cfdfcs.emplace_back(info.funcOp, selectedArchs, numExecs, backwardChannels); } while (!firstCFDFC); return success(); diff --git a/lib/Transforms/BufferPlacement/Utils/CFDFC.cpp b/lib/Transforms/BufferPlacement/Utils/CFDFC.cpp index 2ff2514048..5bb11757bc 100644 --- a/lib/Transforms/BufferPlacement/Utils/CFDFC.cpp +++ b/lib/Transforms/BufferPlacement/Utils/CFDFC.cpp @@ -12,6 +12,7 @@ #include "dynamatic/Transforms/BufferPlacement/Utils/CFDFC.h" #include "dynamatic/Dialect/Handshake/HandshakeOps.h" +#include "dynamatic/Dialect/Handshake/HandshakeTypes.h" #include "dynamatic/Support/CFG.h" #include "dynamatic/Support/ConstraintProgramming/ConstraintProgramming.h" #include "mlir/Dialect/Arith/IR/Arith.h" @@ -20,9 +21,13 @@ #include "mlir/IR/OperationSupport.h" #include "mlir/Support/IndentedOstream.h" #include "mlir/Support/LogicalResult.h" +#include "llvm/ADT/DenseMap.h" +#include "llvm/ADT/DenseSet.h" #include "llvm/ADT/SetOperations.h" #include "llvm/ADT/SmallSet.h" #include "llvm/ADT/SmallVector.h" +#include "llvm/ADT/StringSet.h" +#include "llvm/Support/Casting.h" #include using namespace mlir; @@ -31,6 +36,9 @@ using namespace dynamatic::handshake; using namespace dynamatic::buffer; using namespace dynamatic::experimental; +#include "graphviz/cgraph.h" +#include "graphviz/gvc.h" + namespace { /// Helper data structure to hold mappings between each arch/basic block and the /// Gurobi variable that corresponds to it. @@ -44,6 +52,17 @@ struct MILPVars { }; } // namespace +// AYA: added the following function +bool CFDFC::isCFGCompliant(unsigned srcBB, unsigned dstBB) { + for (size_t i = 0; i < cycle.size(); ++i) { + unsigned nextBB = i == cycle.size() - 1 ? 0 : i + 1; + if (srcBB == cycle[i] && dstBB == cycle[nextBB]) + return true; + } + + return false; +} + /// Initializes all variables in the MILP, one per arch and per basic block. /// Fills in the last argument with mappings between archs/BBs and their /// associated Gurobi variable. @@ -155,87 +174,175 @@ static void setBBConstraints(std::unique_ptr &model, MILPVars &vars) { } } -CFDFC::CFDFC(handshake::FuncOp funcOp, ArchSet &archs, unsigned numExec) +CFDFC::CFDFC(handshake::FuncOp funcOp, ArchSet &archs, unsigned numExec, + DenseSet backwardChannels) : numExecs(numExec) { - // Identify the block that starts the CFDFC; it's the only one that is both - // the source of an arch and the destination of another - std::optional startBB; - llvm::SmallSet uniqueBlocks; - for (ArchBB *arch : archs) { - if (auto [_, inserted] = uniqueBlocks.insert(arch->srcBB); !inserted) { - startBB = arch->srcBB; - break; + bool ayaWay = true; + if (ayaWay) { + llvm::errs() << "\n\tUsing Aya's CFDFC extraction method\n"; + // Identify the block that starts the CFDFC; it's the only one that is both + // the source of an arch and the destination of another + std::optional startBB; + llvm::SmallSet uniqueBlocks; + for (ArchBB *arch : archs) { + if (auto [_, inserted] = uniqueBlocks.insert(arch->srcBB); !inserted) { + startBB = arch->srcBB; + break; + } + if (auto [_, inserted] = uniqueBlocks.insert(arch->dstBB); !inserted) { + startBB = arch->dstBB; + break; + } } - if (auto [_, inserted] = uniqueBlocks.insert(arch->dstBB); !inserted) { - startBB = arch->dstBB; - break; + assert(startBB.has_value() && "failed to identify start of CFDFC"); + + // Form the CFG cycle by stupidly iterating over the archs + cycle.insert(*startBB); + unsigned currentBB = *startBB; + for (size_t i = 0, e = archs.size() - 1; i < e; ++i) { + for (ArchBB *arch : archs) { + if (arch->srcBB == currentBB) { + currentBB = arch->dstBB; + cycle.insert(currentBB); + break; + } + } } - } - assert(startBB.has_value() && "failed to identify start of CFDFC"); + assert(cycle.size() == archs.size() && "failed to construct cycle"); + + // AYA: note to self: the above is common with the old approach + + for (Operation &op : funcOp.getOps()) { + // Get operation's basic block + unsigned srcBB; + if (auto optBB = getLogicBB(&op); !optBB.has_value()) + continue; + else + srcBB = *optBB; + + // The basic block the operation belongs to must be selected in the CFG + // cycle currently under study + if (!cycle.contains(srcBB)) + continue; + + // Add the unit and valid outgoing channels to the CFDFC + units.insert(&op); + for (OpResult res : op.getResults()) { + assert(std::distance(res.getUsers().begin(), res.getUsers().end()) == + 1 && + "value must have unique user"); + + // Get the value's unique user and its basic block + Operation *user = *res.getUsers().begin(); + unsigned dstBB; + if (std::optional optBB = getLogicBB(user); + !optBB.has_value()) + continue; + else + dstBB = *optBB; + + if (!cycle.contains(dstBB)) + continue; + + // AYA: the following if-else is really the core of my logic. First, + // check if it is not a backward edge and the src and dst BBs are part + // of the cycle, consider the channel. Otherwise, if it is a backward + // edge, insert it only if compliant with the CFG + if (!backwardChannels.contains(res)) + channels.insert(res); + else { + // insert backedges only if they are compliant with the CFG + backedges.insert(res); + if (isCFGCompliant(srcBB, dstBB)) + channels.insert(res); + } + } + } + } else { + llvm::errs() << "\n\tUsing the old CFDFC extraction method\n"; - // Form the cycle by stupidly iterating over the archs - cycle.insert(*startBB); - unsigned currentBB = *startBB; - for (size_t i = 0, e = archs.size() - 1; i < e; ++i) { + // Identify the block that starts the CFDFC; it's the only one that is both + // the source of an arch and the destination of another + std::optional startBB; + llvm::SmallSet uniqueBlocks; for (ArchBB *arch : archs) { - if (arch->srcBB == currentBB) { - currentBB = arch->dstBB; - cycle.insert(currentBB); + if (auto [_, inserted] = uniqueBlocks.insert(arch->srcBB); !inserted) { + startBB = arch->srcBB; + break; + } + if (auto [_, inserted] = uniqueBlocks.insert(arch->dstBB); !inserted) { + startBB = arch->dstBB; break; } } - } - assert(cycle.size() == archs.size() && "failed to construct cycle"); - - for (Operation &op : funcOp.getOps()) { - // Get operation's basic block - unsigned srcBB; - if (auto optBB = getLogicBB(&op); !optBB.has_value()) - continue; - else - srcBB = *optBB; - - // The basic block the operation belongs to must be selected - if (!cycle.contains(srcBB)) - continue; - - // Add the unit and valid outgoing channels to the CFDFC - units.insert(&op); - for (OpResult res : op.getResults()) { - assert(std::distance(res.getUsers().begin(), res.getUsers().end()) == 1 && - "value must have unique user"); - - // Get the value's unique user and its basic block - Operation *user = *res.getUsers().begin(); - unsigned dstBB; - if (std::optional optBB = getLogicBB(user); !optBB.has_value()) + assert(startBB.has_value() && "failed to identify start of CFDFC"); + + // Form the cycle by stupidly iterating over the archs + cycle.insert(*startBB); + unsigned currentBB = *startBB; + for (size_t i = 0, e = archs.size() - 1; i < e; ++i) { + for (ArchBB *arch : archs) { + if (arch->srcBB == currentBB) { + currentBB = arch->dstBB; + cycle.insert(currentBB); + break; + } + } + } + assert(cycle.size() == archs.size() && "failed to construct cycle"); + + for (Operation &op : funcOp.getOps()) { + // Get operation's basic block + unsigned srcBB; + if (auto optBB = getLogicBB(&op); !optBB.has_value()) continue; else - dstBB = *optBB; - - if (srcBB != dstBB) { - // The channel is in the CFDFC if it belongs belong to a selected arch - // between two basic blocks - for (size_t i = 0; i < cycle.size(); ++i) { - unsigned nextBB = i == cycle.size() - 1 ? 0 : i + 1; - if (srcBB == cycle[i] && dstBB == cycle[nextBB]) { - channels.insert(res); - if (isCFDFCBackedge(res)) - backedges.insert(res); - break; + srcBB = *optBB; + + // The basic block the operation belongs to must be selected + if (!cycle.contains(srcBB)) + continue; + + // Add the unit and valid outgoing channels to the CFDFC + units.insert(&op); + for (OpResult res : op.getResults()) { + assert(std::distance(res.getUsers().begin(), res.getUsers().end()) == + 1 && + "value must have unique user"); + + // Get the value's unique user and its basic block + Operation *user = *res.getUsers().begin(); + unsigned dstBB; + if (std::optional optBB = getLogicBB(user); + !optBB.has_value()) + continue; + else + dstBB = *optBB; + + if (srcBB != dstBB) { + // The channel is in the CFDFC if it belongs belong to a selected arch + // between two basic blocks + for (size_t i = 0; i < cycle.size(); ++i) { + unsigned nextBB = i == cycle.size() - 1 ? 0 : i + 1; + if (srcBB == cycle[i] && dstBB == cycle[nextBB]) { + channels.insert(res); + if (isCFDFCBackedge(res)) + backedges.insert(res); + break; + } } + } else if (cycle.size() == 1) { + // The channel is in the CFDFC if its producer/consumer belong to the + // same basic block and the CFDFC is just a block looping to itself + channels.insert(res); + if (isCFDFCBackedge(res)) + backedges.insert(res); + } else if (!isBackedge(res)) { + // The channel is in the CFDFC if its producer/consumer belong to the + // same basic block and the channel is not a backedge + channels.insert(res); } - } else if (cycle.size() == 1) { - // The channel is in the CFDFC if its producer/consumer belong to the - // same basic block and the CFDFC is just a block looping to itself - channels.insert(res); - if (isCFDFCBackedge(res)) - backedges.insert(res); - } else if (!isBackedge(res)) { - // The channel is in the CFDFC if its producer/consumer belong to the - // same basic block and the channel is not a backedge - channels.insert(res); } } } @@ -355,6 +462,35 @@ void dynamatic::buffer::getDisjointBlockUnions( } } +// AYA: Added this from Jiahui to write the CFDFCs in DOT for debugging +void CFDFC::writeDot(const std::string &fileName) { + Agraph_t *gv = agopen(const_cast("cfdfc"), Agdirected, nullptr); + for (auto value : channels) { + auto *pred = value.getDefiningOp(); + auto succ = value.getUsers().begin(); + std::string predName = + pred->getAttrOfType(NameAnalysis::ATTR_NAME).str(); + std::string succName = + succ->getAttrOfType(NameAnalysis::ATTR_NAME).str(); + + auto *predNode = agnode(gv, const_cast(predName.c_str()), + /* create if not exist */ 1); + + auto *succNode = agnode(gv, const_cast(succName.c_str()), + /* create if not exist */ 1); + + agedge(gv, predNode, succNode, nullptr, 1); + } + // Write to DOT file + FILE *fs = fopen(fileName.c_str(), "w"); + + if (!fs) + llvm::errs() << "Failed to write file\n"; + + agwrite(gv, fs); + fclose(fs); +} + LogicalResult dynamatic::buffer::extractCFDFC( handshake::FuncOp funcOp, ArchSet &archs, BBSet &bbs, ArchSet &selectedArchs, unsigned &numExecs, CPSolver::SolverKind solverKind, diff --git a/tools/unit-generators/vhdl/generators/handshake/init.py b/tools/unit-generators/vhdl/generators/handshake/init.py new file mode 100644 index 0000000000..256c2e988c --- /dev/null +++ b/tools/unit-generators/vhdl/generators/handshake/init.py @@ -0,0 +1,157 @@ +from generators.support.signal_manager import generate_concat_signal_manager +from generators.support.signal_manager.utils.concat import get_concat_extra_signals_bitwidth + + +def generate_init(name, params): + bitwidth = params["bitwidth"] + extra_signals = params.get("extra_signals", None) + # The initial value to use for the buffer + initial_value = params.get("initial_value", 0) + + if extra_signals: + return _generate_init_signal_manager(name, bitwidth, extra_signals, initial_value) + elif bitwidth == 0: + return _generate_init_dataless(name) + else: + return _generate_init(name, bitwidth, initial_value) + + +def _generate_init_dataless(name): + + entity = f""" +library ieee; +use ieee.std_logic_1164.all; +use ieee.numeric_std.all; + +-- Entity of init_dataless +entity {name} is + port ( + clk, rst : in std_logic; + -- input channel + ins_valid : in std_logic; + ins_ready : out std_logic; + -- output channel + outs_valid : out std_logic; + outs_ready : in std_logic + ); +end entity; +""" + + architecture = f""" +-- Architecture of init_dataless +architecture arch of {name} is + signal fullReg, outputValid : std_logic; +begin + outputValid <= ins_valid or fullReg; + + process (clk) is + begin + if (rising_edge(clk)) then + if (rst = '1') then + fullReg <= '1'; + else + fullReg <= outputValid and not outs_ready; + end if; + end if; + end process; + + ins_ready <= not fullReg; + outs_valid <= outputValid; +end architecture; +""" + + return entity + architecture + + +def _generate_init(name, bitwidth, initial_value): + init_dataless_name = f"{name}_dataless" + + dependencies = _generate_init_dataless( + init_dataless_name) + + dataReg_init = f"'{initial_value}'" + + entity = f""" +library ieee; +use ieee.std_logic_1164.all; +use ieee.numeric_std.all; + +-- Entity of init +entity {name} is + port ( + clk, rst : in std_logic; + -- input channel + ins : in std_logic_vector({bitwidth} - 1 downto 0); + ins_valid : in std_logic; + ins_ready : out std_logic; + -- output channel + outs : out std_logic_vector({bitwidth} - 1 downto 0); + outs_valid : out std_logic; + outs_ready : in std_logic + ); +end entity; +""" + + architecture = f""" +-- Architecture of init +architecture arch of {name} is + signal regEnable, regNotFull : std_logic; + signal dataReg : std_logic_vector({bitwidth} - 1 downto 0); +begin + regEnable <= regNotFull and ins_valid and not outs_ready; + + control : entity work.{init_dataless_name} + port map( + clk => clk, + rst => rst, + ins_valid => ins_valid, + ins_ready => regNotFull, + outs_valid => outs_valid, + outs_ready => outs_ready + ); + + process (clk) is + begin + if (rising_edge(clk)) then + if (rst = '1') then + dataReg <= (others => {dataReg_init}); + elsif (regEnable) then + dataReg <= ins; + end if; + end if; + end process; + + process (regNotFull, dataReg, ins) is + begin + if (regNotFull) then + outs <= ins; + else + outs <= dataReg; + end if; + end process; + + ins_ready <= regNotFull; + +end architecture; +""" + + return dependencies + entity + architecture + + +def _generate_init_signal_manager(name, bitwidth, extra_signals, initial_value): + extra_signals_bitwidth = get_concat_extra_signals_bitwidth(extra_signals) + return generate_concat_signal_manager( + name, + [{ + "name": "ins", + "bitwidth": bitwidth, + "extra_signals": extra_signals + }], + [{ + "name": "outs", + "bitwidth": bitwidth, + "extra_signals": extra_signals + }], + extra_signals, + lambda name: _generate_init(name, bitwidth + extra_signals_bitwidth, initial_value)) + \ No newline at end of file diff --git a/tools/unit-generators/vhdl/vhdl-unit-generator.py b/tools/unit-generators/vhdl/vhdl-unit-generator.py index 0d1ee570f0..a7ad974e75 100644 --- a/tools/unit-generators/vhdl/vhdl-unit-generator.py +++ b/tools/unit-generators/vhdl/vhdl-unit-generator.py @@ -136,5 +136,6 @@ def main(generators): generators.add("handshake", "remsi") generators.add("handshake", "ram") generators.add("handshake", "sharing_wrapper") + generators.add("handshake", "init") main(generators)