From cc1e03a0af5c6b16ec18c1099c286727f8bd561a Mon Sep 17 00:00:00 2001 From: 1801ZDL <1241272204@qq.com> Date: Mon, 10 Aug 2026 10:43:09 +0800 Subject: [PATCH 1/9] feat: support scalar V->C dependencies in dynamic CV pipeline Transfer tensor.extract scalar results from VECTOR blocks to CUBE blocks via SSBuffer with PIPE_S sync, so CUBE blocks no longer recompute VECTOR-only math chains (math.floor/math.ceil) or depend on them across the CV boundary. - DataDependencyAnalysis: detect scalar V->C deps from CUBE-side tensor.extract of VECTOR-only math ops (using the op's own core_type, since a block may mix VECTOR and CUBE ops), tracing the downstream pure scalar chain (fptosi/muli/subi/addi) to find CUBE consumers, including uses inside for/if bodies. - InterCoreTransferAndSync: insert SSBuffer store/load pairs with cross-core deps and PIPE_S sync for scalar transfers. The store is placed right after the extract, the load after the store when they share a block, and the store is skipped in the consumer-use rewrite so it keeps referencing the original extract (no store->load self loop). - SeparateCVScope: keep VECTOR-only boundary chains out of the CUBE scope by rewriting leftover VECTOR references to the CUBE-side load chains and retaining only ops each scope actually needs. - OpClassifier: classify math.floor/ceil as VECTOR; isVectorOnlyOp covers math::CeilOp. - SSBufferManager: the SSBuffer memref element type must match the stored value (memref.store verifies value/element types); the LLVM pointer era allowed a hardcoded i32 memref. Rebased onto latest main-dev: adapted LLVM::LoadOp/StoreOp to memref::LoadOp/StoreOp (the main-dev LLVM->memref migration) and dropped the now-removed DependencyInfo::isScaler flag. Co-Authored-By: DeLong code --- .../Common/SSBufferManager.h | 16 +- .../SplitDataflow/DataDependencyAnalysis.h | 1 + .../UpdateConditionInfo.cpp | 5 +- .../Common/SSBufferManager.cpp | 9 +- .../lib/DynamicCVPipeline/Common/Utils.cpp | 2 +- .../PlanComputeBlock/OpClassifier.cpp | 40 +- .../SplitDataflow/DataDependencyAnalysis.cpp | 364 +++++++++++++++++- .../InterCoreTransferAndSync.cpp | 38 +- .../SplitDataflow/SeparateCVScope.cpp | 292 ++++++++++++++ 9 files changed, 733 insertions(+), 34 deletions(-) diff --git a/third_party/ascend/include/DynamicCVPipeline/Common/SSBufferManager.h b/third_party/ascend/include/DynamicCVPipeline/Common/SSBufferManager.h index 855124da54..b9a49424a2 100644 --- a/third_party/ascend/include/DynamicCVPipeline/Common/SSBufferManager.h +++ b/third_party/ascend/include/DynamicCVPipeline/Common/SSBufferManager.h @@ -90,27 +90,31 @@ class SSBufferManager { static constexpr int ADDR_INT_TYPE = 64; static constexpr int CONST_INT_TYPE = 32; -inline MemRefType getSsbufMemrefType(Builder &builder) { - auto i32Type = builder.getIntegerType(CONST_INT_TYPE); +inline MemRefType getSsbufMemrefType(Builder &builder, Type elemType) { auto addressSpaceAttr = builder.getAttr(hivm::AddressSpace::SSBUF); - return MemRefType::get({}, i32Type, nullptr, addressSpaceAttr); + return MemRefType::get({}, elemType, nullptr, addressSpaceAttr); } inline std::pair -getSsbufConstAndPointerCast(OpBuilder &builder, Location loc, uint64_t addr) { +getSsbufConstAndPointerCast(OpBuilder &builder, Location loc, uint64_t addr, + Type elemType) { auto i64Type = builder.getIntegerType(ADDR_INT_TYPE); auto addrAttr = builder.getIntegerAttr(i64Type, addr); auto addrConst = builder.create(loc, i64Type, addrAttr); return {addrConst, - builder.create(loc, getSsbufMemrefType(builder), + builder.create(loc, + getSsbufMemrefType(builder, + elemType), addrConst.getResult())}; } inline hivm::PointerCastOp createPointerCastOp(OpBuilder &builder, Location loc, uint64_t addr) { - return getSsbufConstAndPointerCast(builder, loc, addr).second; + // Default to i32 for callers that only store i32 values into SSBuffer. + auto i32Type = builder.getIntegerType(CONST_INT_TYPE); + return getSsbufConstAndPointerCast(builder, loc, addr, i32Type).second; } } // namespace triton diff --git a/third_party/ascend/include/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.h b/third_party/ascend/include/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.h index 15b120df57..b2392219df 100644 --- a/third_party/ascend/include/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.h +++ b/third_party/ascend/include/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.h @@ -144,6 +144,7 @@ class DataDependencyAnalysisPass mlir::Operation *predOp, mlir::Operation *nextOp); void analyzeExternalInputs(DataDependencyInfo &info); void analyzeExternalOutputs(DataDependencyInfo &info); + void analyzeScalarVToCDependencies(DataDependencyInfo &info); void analyzeMemoryEffect(DataDependencyInfo &info); std::pair findCommonLevelBlockIds(DataDependencyInfo &info, diff --git a/third_party/ascend/lib/DynamicCVPipeline/AddControlFlowCondition/UpdateConditionInfo.cpp b/third_party/ascend/lib/DynamicCVPipeline/AddControlFlowCondition/UpdateConditionInfo.cpp index 03cbd4aef7..70e5ad24b8 100644 --- a/third_party/ascend/lib/DynamicCVPipeline/AddControlFlowCondition/UpdateConditionInfo.cpp +++ b/third_party/ascend/lib/DynamicCVPipeline/AddControlFlowCondition/UpdateConditionInfo.cpp @@ -95,7 +95,7 @@ UpdateConditionInfoPass::allocSSBuffer(ModuleOp module) { OpBuilder builder(module.getContext()); auto i64Type = builder.getIntegerType(ADDR_INT_TYPE); auto i32Type = builder.getIntegerType(CONST_INT_TYPE); - auto memrefType = getSsbufMemrefType(builder); + auto memrefType = getSsbufMemrefType(builder, i32Type); // alloc 2 group of ssbuffer pointers: // Core Vector 0: allocate ssbuffer address: 0, 4, 8, ... @@ -489,7 +489,8 @@ UpdateConditionInfoPass::computeVectorSSBufferMemrefs( auto ssbAddr = builder.create(loc, ssbBaseAddr, ssbAddrOffset); Value memref = builder.create( - loc, getSsbufMemrefType(builder), ssbAddr.getResult()); + loc, getSsbufMemrefType(builder, builder.getIntegerType(CONST_INT_TYPE)), + ssbAddr.getResult()); vectorSSBufferMemrefs[groupIdx] = memref; } diff --git a/third_party/ascend/lib/DynamicCVPipeline/Common/SSBufferManager.cpp b/third_party/ascend/lib/DynamicCVPipeline/Common/SSBufferManager.cpp index 5b7e76aa02..4a4761f8a5 100644 --- a/third_party/ascend/lib/DynamicCVPipeline/Common/SSBufferManager.cpp +++ b/third_party/ascend/lib/DynamicCVPipeline/Common/SSBufferManager.cpp @@ -103,8 +103,11 @@ SSBufferManager::writeToSSBuffer(Value value, OpBuilder &builder, int64_t addrValue = addrResult.value(); Location loc = builder.getUnknownLoc(); + // The memref element type must match the stored value's type: with LLVM + // pointer ops the element type was untyped, but memref.store verifies that + // the value type matches the memref element type. auto [constOp, pointerCastOp] = - getSsbufConstAndPointerCast(builder, loc, addrValue); + getSsbufConstAndPointerCast(builder, loc, addrValue, value.getType()); createdOps.push_back(constOp); createdOps.push_back(pointerCastOp); @@ -129,8 +132,8 @@ SSBufferManager::readFromSSBuffer(int64_t addr, OpBuilder &builder, } Location loc = builder.getUnknownLoc(); - auto [constOp, pointerCastOp] = - getSsbufConstAndPointerCast(builder, loc, addr); + auto [constOp, pointerCastOp] = getSsbufConstAndPointerCast( + builder, loc, addr, findResult.value().second); createdOps.push_back(constOp); createdOps.push_back(pointerCastOp); diff --git a/third_party/ascend/lib/DynamicCVPipeline/Common/Utils.cpp b/third_party/ascend/lib/DynamicCVPipeline/Common/Utils.cpp index 37959eeb2f..e13886992e 100644 --- a/third_party/ascend/lib/DynamicCVPipeline/Common/Utils.cpp +++ b/third_party/ascend/lib/DynamicCVPipeline/Common/Utils.cpp @@ -108,7 +108,7 @@ bool isVectorOnlyOp(Operation *op) { return llvm::TypeSwitch(op) .Case([](linalg::ReduceOp) { return true; }) - .Case([](Operation *op) { + .Case([](Operation *op) { return isa(op->getResult(0).getType()); }) .Default([](auto) { return false; }); diff --git a/third_party/ascend/lib/DynamicCVPipeline/PlanComputeBlock/OpClassifier.cpp b/third_party/ascend/lib/DynamicCVPipeline/PlanComputeBlock/OpClassifier.cpp index f2043fb625..34e1e18ecb 100644 --- a/third_party/ascend/lib/DynamicCVPipeline/PlanComputeBlock/OpClassifier.cpp +++ b/third_party/ascend/lib/DynamicCVPipeline/PlanComputeBlock/OpClassifier.cpp @@ -703,6 +703,20 @@ void OpClassifierPass::getUpstreamOpsWithMemoryDeps( } } +// An arith/math op with a tensor result is VECTOR-only and must not be marked +// CUBE; scalar arith/math (index/i32 computation) may still be marked CUBE. +static bool isTensorArithOrMathOp(Operation *op) { + if (!isa(op->getDialect())) { + return false; + } + for (Value result : op->getResults()) { + if (isa(result.getType())) { + return true; + } + } + return false; +} + // Propagate CUBE core type upstream int OpClassifierPass::propagateCubeUpstream() { LLVM_DEBUG(DBGS() << "--- Step 2: CUBE upstream BFS --->\n"); @@ -729,21 +743,12 @@ int OpClassifierPass::propagateCubeUpstream() { if (!def || cubeVisited.count(def) || isa(def)) continue; - // Skip arith dialect ops with tensor results (they should be VECTOR, not - // CUBE) - if (isa(def->getDialect())) { - bool hasTensorResult = false; - for (Value result : def->getResults()) { - if (isa(result.getType())) { - hasTensorResult = true; - break; - } - } - if (hasTensorResult) { - LLVM_DEBUG(DBGS() << "skip " << def->getName().getStringRef() - << ": arith tensor op\n"); - continue; - } + // Skip arith/math ops with tensor results (they are VECTOR-only, not + // CUBE); scalar arith/math may still be marked CUBE. + if (isTensorArithOrMathOp(def)) { + LLVM_DEBUG(DBGS() << "skip " << def->getName().getStringRef() + << ": arith/math tensor op\n"); + continue; } // Skip operations inside linalg block (internal values) @@ -927,6 +932,11 @@ void OpClassifierPass::propagateCubeUpstreamForOp(Operation *startOp) { continue; if (isa(upstreamOp)) continue; + // Align with the main propagateCubeUpstream: only skip arith/math ops + // with tensor results (they are VECTOR-only); scalar arith/math ops + // (index/i32 computation) may still be marked CUBE. + if (isTensorArithOrMathOp(upstreamOp)) + continue; cubeVisited.insert(upstreamOp); LLVM_DEBUG(DBGS() << "\t\tcube upstream: " diff --git a/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.cpp b/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.cpp index 74aacd3a6d..3f55880e1c 100644 --- a/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.cpp +++ b/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.cpp @@ -163,7 +163,12 @@ bool DataDependencyAnalysisPass::isValid1DValueForDependency( mlir::Value value) { auto tensorTy = dyn_cast(value.getType()); if (tensorTy && tensorTy.getRank() == SHAPE_1D_LENGTH) { - return true; + // Only 1-D tensors consumed by linalg.broadcast count as valid cross-core + // dependencies. A 1-D tensor consumed by tensor.extract is scalarized + // first and passed through the scalar SSBuffer dependency channel, so it + // does not need a separate 1-D tensor CopyOp. + return llvm::all_of(value.getUsers(), + [](mlir::Operation *u) { return isa(u); }); } return false; } @@ -841,6 +846,361 @@ void DataDependencyAnalysisPass::collectMemDepInfo( memoryDependencies.push_back(depInfo); } +// Trace defining op chain to check if any upstream op is vector-only. +// Returns true if a vector-only op is found in the chain. +// Values in `stopValues` are treated as already-handled scalar dependencies; +// the trace stops when it hits one, avoiding redundant detection of downstream +// scalars whose upstream values have already been transferred via SSBuffer. +// This mirrors AnalyzeCubeControlFlowInputChain's hasIncompatibleOpForCondition. +static bool hasVectorOpInDefChain( + mlir::Value val, + llvm::DenseSet &visited, + const llvm::DenseSet *stopValues = nullptr) { + if (stopValues && stopValues->contains(val)) { + return false; + } + + mlir::Operation *defOp = val.getDefiningOp(); + if (!defOp || visited.contains(defOp)) { + return false; + } + visited.insert(defOp); + + if (CVPipeline::isVectorOnlyOp(defOp)) { + return true; + } + + for (mlir::Value operand : defOp->getOperands()) { + if (hasVectorOpInDefChain(operand, visited, stopValues)) { + return true; + } + } + return false; +} + +// Check if a scf.for body contains CUBE operations. +static bool forOpHasCubeOps( + scf::ForOp forOp, + llvm::DenseMap &blockInfoMap) { + bool hasCube = false; + forOp.walk([&](mlir::Operation *op) { + auto blockIdOpt = CVPipeline::getOpBlockId(op); + if (!blockIdOpt) { + return mlir::WalkResult::advance(); + } + auto it = blockInfoMap.find(*blockIdOpt); + if (it != blockInfoMap.end() && it->second.isCube) { + hasCube = true; + return mlir::WalkResult::interrupt(); + } + return mlir::WalkResult::advance(); + }); + return hasCube; +} + +// Check if any region of an scf.if contains CUBE operations. +static bool ifOpHasCubeOps( + scf::IfOp ifOp, + llvm::DenseMap &blockInfoMap) { + bool hasCube = false; + ifOp.walk([&](mlir::Operation *op) { + auto blockIdOpt = CVPipeline::getOpBlockId(op); + if (!blockIdOpt) { + return mlir::WalkResult::advance(); + } + auto it = blockInfoMap.find(*blockIdOpt); + if (it != blockInfoMap.end() && it->second.isCube) { + hasCube = true; + return mlir::WalkResult::interrupt(); + } + return mlir::WalkResult::advance(); + }); + return hasCube; +} + +// Analyze scalar V->C dependencies from control flow ops. +// Detects when scf.for loop bounds or scf.if conditions are scalar values whose +// defining chain traces back to vector-only ops (e.g. math.floor/math.ceil on +// tensors, linalg.reduce). These scalars must be transferred from VECTOR to CUBE. +// +// To avoid redundant transfers, scalars already recorded as dependencies are +// tracked in a stop-set: downstream scalars whose defining chain only reaches +// vector-only ops through an already-handled scalar are NOT re-recorded. +void DataDependencyAnalysisPass::analyzeScalarVToCDependencies( + DataDependencyInfo &info) { + auto &blockInfoMap = info.getBlockInfoMap(); + auto &v2cDependencies = info.getV2CDependencies(); + + // Scalars already recorded as V->C deps — downstream values depending on + // these don't need separate transfer since the upstream value will be + // available on the CUBE side after SSBuffer transfer. + llvm::DenseSet handledScalarValues; + + // For-loops whose bounds have been handled — ifOp conditions inside these + // loops will be computable on CUBE via the induction variable after the + // bounds are transferred, so they don't need separate scalar deps. + llvm::DenseSet handledForOps; + + LOG_DEBUG("Analyzing scalar V->C dependencies from control flow ops...\n"); + + // ---- tensor.extract ops ---- + // Detect scalars extracted from VECTOR-produced tensors and consumed by CUBE + // blocks. The extract result is the natural scalar dependency boundary: + // transferring it via SSBuffer avoids the need for 1-D tensor CopyOps. + module.walk([&](tensor::ExtractOp extractOp) { + mlir::Value sourceTensor = extractOp.getTensor(); + mlir::Operation *tensorDefOp = sourceTensor.getDefiningOp(); + if (!tensorDefOp) { + return; + } + + // Only handle extracts whose source tensor is produced in a VECTOR block. + auto tensorBlockIdOpt = CVPipeline::getOpBlockId(tensorDefOp); + if (!tensorBlockIdOpt) { + return; + } + auto tensorBlockIt = blockInfoMap.find(*tensorBlockIdOpt); + if (tensorBlockIt == blockInfoMap.end() || tensorBlockIt->second.isCube) { + return; + } + // tensorDefOp must be a VECTOR-only op on tensor (e.g. math.floor/ceil). + if (!CVPipeline::isVectorOnlyOp(tensorDefOp)) { + return; + } + + // Only handle extracts located in a CUBE block. The VECTOR-side extract + // (e.g. block 21) is the original computation; the CUBE-side extract + // (e.g. block 15) is the copy whose scalar result the CUBE side actually + // consumes. Transferring the VECTOR-side extract creates a CUBE load with + // no users (redundant transfer). + auto extractBlockIdOpt = CVPipeline::getOpBlockId(extractOp.getOperation()); + if (!extractBlockIdOpt) { + return; + } + auto extractBlockIt = blockInfoMap.find(*extractBlockIdOpt); + if (extractBlockIt == blockInfoMap.end() || + !extractBlockIt->second.isCube) { + return; + } + + mlir::Value scalarResult = extractOp.getResult(); + if (!isa(scalarResult.getType())) { + return; + } + + // The scalar must be consumed (directly or through a downstream pure + // scalar chain, e.g. fptosi/muli/subi producing loop bounds) in at least + // one CUBE block or in a for/if with CUBE content. + int producerId = *extractBlockIdOpt; + + llvm::DenseSet handledConsumers; + bool hasCubConsumer = false; + llvm::SmallVector worklist; + llvm::DenseSet visited; + worklist.push_back(scalarResult); + while (!worklist.empty()) { + mlir::Value cur = worklist.pop_back_val(); + if (!visited.insert(cur).second) { + continue; + } + for (mlir::Operation *user : cur.getUsers()) { + auto userBlockIdOpt = CVPipeline::getOpBlockId(user); + bool usedInCube = false; + if (userBlockIdOpt) { + auto it = blockInfoMap.find(*userBlockIdOpt); + if (it != blockInfoMap.end() && it->second.isCube) { + usedInCube = true; + handledConsumers.insert(*userBlockIdOpt); + } + } + if (!usedInCube) { + // A for/if containing CUBE ops also counts as a CUBE consumer. + if (auto forOp = dyn_cast(user)) { + if (forOpHasCubeOps(forOp, blockInfoMap) && userBlockIdOpt) { + usedInCube = true; + handledConsumers.insert(*userBlockIdOpt); + } + } else if (auto ifOp = dyn_cast(user)) { + if (ifOpHasCubeOps(ifOp, blockInfoMap) && userBlockIdOpt) { + usedInCube = true; + handledConsumers.insert(*userBlockIdOpt); + } + } + } + if (usedInCube) { + hasCubConsumer = true; + continue; + } + // Follow pure scalar compute chain (single-result, no regions). + if (user->getNumRegions() == 0 && user->getNumResults() == 1 && + user->getResult(0).getType().isIntOrIndexOrFloat()) { + worklist.push_back(user->getResult(0)); + } + } + } + if (hasCubConsumer) { + // Create a single dependency for this extract result. All consumers + // share one SSBuffer store; one dep avoids duplicate stores (and the + // duplicate sync that can deadlock the cores). + if (!handledConsumers.empty()) { + int consumerId = *handledConsumers.begin(); + collectDepInfo(scalarResult, DependencyType::VectorToCube, + v2cDependencies, producerId, consumerId, info); + } + LOG_DEBUG("Found scalar V->C dependency from tensor.extract: " + << scalarResult << "\n"); + handledScalarValues.insert(scalarResult); + // Mark any enclosing for-loop as handled: its bounds and any ifOp + // conditions inside derive from this transferred scalar, so they don't + // need separate transfers. + mlir::Operation *parent = extractOp->getParentOp(); + while (parent) { + if (auto parentForOp = dyn_cast(parent)) { + handledForOps.insert(parentForOp); + break; + } + parent = parent->getParentOp(); + } + } + }); + + // ---- scf.for loop bounds ---- + module.walk([&](scf::ForOp forOp) { + if (!forOpHasCubeOps(forOp, blockInfoMap)) { + return; + } + + llvm::SmallVector bounds; + bounds.push_back(forOp.getLowerBound()); + bounds.push_back(forOp.getUpperBound()); + // Step is typically a constant; still check it for completeness. + bounds.push_back(forOp.getStep()); + + for (mlir::Value bound : bounds) { + // Only handle scalar types (int/float/index). Tensor types cannot be + // transferred through the SSBuffer scalar channel. + if (!isa( + bound.getType())) { + continue; + } + + mlir::Operation *defOp = bound.getDefiningOp(); + if (!defOp) { + continue; + } + + llvm::DenseSet visited; + bool hasVectorDep = + hasVectorOpInDefChain(bound, visited, &handledScalarValues); + if (!hasVectorDep) { + // The trace may have been stopped by an already-handled scalar + // (e.g. a tensor.extract result transferred by the extract pass). + // In that case the bound itself needs no new transfer, but the + // enclosing loop still becomes "handled" so that ifOp conditions + // inside it are not redundantly detected. + llvm::DenseSet visitedNoStop; + if (!hasVectorOpInDefChain(bound, visitedNoStop)) { + continue; + } + handledForOps.insert(forOp); + continue; + } + + auto producerIdOpt = CVPipeline::getOpBlockId(defOp); + if (!producerIdOpt) { + continue; + } + int producerId = *producerIdOpt; + + LOG_DEBUG("Found scalar V->C dependency from forOp bounds: " + << bound << "\n"); + + // Mark this for-loop as handled so that ifOp conditions nested inside + // it are skipped — they will be computable on CUBE via the induction + // variable once the bounds are transferred. + handledForOps.insert(forOp); + + // Record a single V->C dependency using the for-loop op's own block_id + // as the consumer. The bound is also used by ops inside the loop body, + // but those are in the same CUBE scope after CV separation and will use + // the transferred value via normal SSA dominance — no separate transfer + // needed per inner consumer. + auto forBlockIdOpt = CVPipeline::getOpBlockId(forOp.getOperation()); + if (!forBlockIdOpt) { + continue; + } + int forBlockId = *forBlockIdOpt; + auto it = blockInfoMap.find(forBlockId); + if (it == blockInfoMap.end() || !it->second.isCube) { + continue; + } + collectDepInfo(bound, DependencyType::VectorToCube, v2cDependencies, + producerId, forBlockId, info); + handledScalarValues.insert(bound); + } + }); + + // ---- scf.if condition ---- + module.walk([&](scf::IfOp ifOp) { + if (!ifOpHasCubeOps(ifOp, blockInfoMap)) { + return; + } + + // Skip if this ifOp is nested inside a for-loop whose bounds have already + // been transferred. The induction variable carries the transferred values, + // so any scalar condition computed from it is already available on CUBE. + mlir::Operation *parent = ifOp->getParentOp(); + while (parent) { + if (auto parentForOp = dyn_cast(parent)) { + if (handledForOps.contains(parentForOp)) { + return; + } + } + parent = parent->getParentOp(); + } + + mlir::Value condition = ifOp.getCondition(); + if (!isa(condition.getType())) { + return; + } + + mlir::Operation *defOp = condition.getDefiningOp(); + if (!defOp) { + return; + } + + llvm::DenseSet visited; + if (!hasVectorOpInDefChain(condition, visited, &handledScalarValues)) { + return; + } + + auto producerIdOpt = CVPipeline::getOpBlockId(defOp); + if (!producerIdOpt) { + return; + } + int producerId = *producerIdOpt; + + auto consumerIdOpt = CVPipeline::getOpBlockId(ifOp.getOperation()); + if (!consumerIdOpt) { + return; + } + int consumerId = *consumerIdOpt; + auto it = blockInfoMap.find(consumerId); + if (it == blockInfoMap.end() || !it->second.isCube) { + return; + } + + LOG_DEBUG("Found scalar V->C dependency from ifOp condition: " + << condition << "\n"); + + collectDepInfo(condition, DependencyType::VectorToCube, v2cDependencies, + producerId, consumerId, info); + handledScalarValues.insert(condition); + }); + + LOG_DEBUG("Scalar V->C dependency analysis complete.\n"); +} + void DataDependencyAnalysisPass::analyzeMemoryEffect(DataDependencyInfo &info) { auto &memoryDependencies = info.getMemoryDependencies(); LOG_DEBUG("\n=== start mem dep analysis ===\n"); @@ -1055,6 +1415,8 @@ void DataDependencyAnalysisPass::runOnOperation() { analyzeExternalOutputs(info); + analyzeScalarVToCDependencies(info); + // Step 4: Analyze memory dependencies (memdep sync) analyzeMemoryEffect(info); diff --git a/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/InterCoreTransferAndSync.cpp b/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/InterCoreTransferAndSync.cpp index 36ef8df735..db4f9d2266 100644 --- a/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/InterCoreTransferAndSync.cpp +++ b/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/InterCoreTransferAndSync.cpp @@ -547,7 +547,15 @@ Operation *InterCoreTransferAndSyncPass::insertVectorToCubeTransfer( int cubeBlockId = CVPipeline::getOpBlockId(cubeStartOp).value_or(-1); if (isScalarDependency(dep.value)) { - builder.setInsertionPointAfter(vectorEndOp); + // Insert the store right after srcValue's defining op (when it has one) + // so that the store dominates the load and the scalar consumers. Falling + // back to vectorEndOp keeps behavior for block-argument / external values. + mlir::Operation *srcDefOp = srcValue.getDefiningOp(); + if (srcDefOp) { + builder.setInsertionPointAfter(srcDefOp); + } else { + builder.setInsertionPointAfter(vectorEndOp); + } SmallVector writeOps; LOG_DEBUG("before writeToSSBuffer\n"); auto addrOpt = ssbufferManager.writeToSSBuffer(srcValue, builder, writeOps); @@ -568,7 +576,16 @@ Operation *InterCoreTransferAndSyncPass::insertVectorToCubeTransfer( attachCrossCoreDeps(sendOp, transferIndex, CVPipeline::crossCoreProducerId, builder); LOG_DEBUG("before readFromSSBuffer\n"); - builder.setInsertionPoint(cubeStartOp); + // Place the load after the store when both are in the same MLIR block, + // otherwise the load would read an uninitialized SSBuffer slot (store and + // load can share a block when producer block == consumer block). + if (sendOp && cubeStartOp && + sendOp->getBlock() == cubeStartOp->getBlock() && + !sendOp->isBeforeInBlock(cubeStartOp)) { + builder.setInsertionPointAfter(sendOp); + } else { + builder.setInsertionPoint(cubeStartOp); + } SmallVector readOps; auto loadedValueOpt = ssbufferManager.readFromSSBuffer(addr, builder, readOps); @@ -655,6 +672,12 @@ Operation *InterCoreTransferAndSyncPass::insertVectorToCubeTransfer( } for (Operation *user : users) { LOG_DEBUG("[v->c user]" << *user << "\n"); + // Do not rewrite the store/send op itself: it must keep referencing the + // original srcValue, otherwise the value stored into SSBuffer would be the + // just-loaded receiveValue (a store→load self loop). + if (user == sendOp) { + continue; + } auto userBlockIdOpt = CVPipeline::getOpBlockId(user); if (userBlockIdOpt && *userBlockIdOpt == dep.iniConsumerBlockId) { user->replaceUsesOfWith(srcValue, receiveValue); @@ -789,10 +812,13 @@ InterCoreTransferAndSyncPass::getTransferPipeConfig(Operation *transferOp, config.srcCoreType = "VECTOR"; config.dstCoreType = "CUBE"; } else if (isa(transferOp)) { - config.forReadTPipe = pipeVAttr; - config.forReadPipe = pipeFixAttr; - config.forWriteTPipe = pipeFixAttr; - config.forWritePipe = pipeVAttr; + // Scalar transfers use PIPE_S (scalar pipe). Using PIPE_V/PIPE_FIX for + // these syncs shares flag space with vector/fix tensor transfers and can + // deadlock the cores; PIPE_S keeps the scalar sync isolated. + config.forReadTPipe = pipeSAttr; + config.forReadPipe = pipeSAttr; + config.forWriteTPipe = pipeSAttr; + config.forWritePipe = pipeSAttr; config.srcCoreAttr = vecCoreAttr; config.dstCoreAttr = cubeCoreAttr; config.srcCoreType = "VECTOR"; diff --git a/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/SeparateCVScope.cpp b/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/SeparateCVScope.cpp index 46505d853c..cd61c417cb 100644 --- a/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/SeparateCVScope.cpp +++ b/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/SeparateCVScope.cpp @@ -22,6 +22,8 @@ #include +#include + #include "llvm/ADT/SmallPtrSet.h" #include "llvm/ADT/SmallVector.h" #include "llvm/Support/Debug.h" @@ -31,13 +33,16 @@ #include "bishengir/Dialect/Scope/IR/Scope.h" #include "mlir/Dialect/Arith/IR/Arith.h" #include "mlir/Dialect/Func/IR/FuncOps.h" +#include "mlir/Dialect/LLVMIR/LLVMDialect.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" #include "mlir/Dialect/SCF/IR/SCF.h" +#include "mlir/Dialect/Tensor/IR/Tensor.h" #include "mlir/IR/Builders.h" #include "mlir/IR/BuiltinAttributes.h" #include "mlir/IR/IRMapping.h" #include "mlir/Interfaces/LoopLikeInterface.h" #include "mlir/Pass/Pass.h" +#include "mlir/Transforms/RegionUtils.h" #include "ascend/include/DynamicCVPipeline/Common/Utils.h" #include "ascend/include/DynamicCVPipeline/SplitDataflow/SeparateCVScope.h" @@ -951,6 +956,234 @@ static void cleanupSsbufferAttrs(Operation *rootOp) { rootOp->walk([](Operation *op) { removeSsbufferAttrs(op); }); } +// Record, for every value produced along the forward use-chain rooted at +// `root`, the sequence of op names that led to it. This lets us pair up a +// value on the VECTOR side (from a SSBuffer store) with the structurally +// identical value on the CUBE side (from the matching SSBuffer load). +static void collectChainPaths( + Value root, llvm::DenseMap &paths) { + llvm::SmallVector worklist; + worklist.push_back(root); + paths[root] = ""; + while (!worklist.empty()) { + Value cur = worklist.pop_back_val(); + std::string curPath = paths[cur]; + Region *rootRegion = root.getParentRegion(); + for (Operation *user : cur.getUsers()) { + if (user->getNumResults() == 0) { + continue; + } + // Only follow single-result pure-compute ops; stop at anything with + // regions or side effects. + if (user->getNumRegions() > 0 || + !mlir::wouldOpBeTriviallyDead(user)) { + continue; + } + // Do not cross into nested regions (e.g. a scf.for body): a use inside a + // loop is not a candidate for the boundary-mapping used here, and + // substituting a value defined there would break dominance. + if (user->getBlock()->getParent() != rootRegion) { + continue; + } + Value res = user->getResult(0); + if (paths.contains(res)) { + continue; + } + paths[res] = curPath + user->getName().getStringRef().str() + ";"; + worklist.push_back(res); + } + } +} + +// Rewrite uses in the CUBE scope that reference values computed on the VECTOR +// side (boundary scalars derived from math.floor/math.ceil via tensor.extract). +// +// Such references are left-over copies of the VECTOR boundary chain that +// SeparateCVScope cloned into the CUBE scope's mixed for-loop. The CUBE scope +// already has an equivalent chain rooted at the SSBuffer load for the same +// transfer_id (e.g. load %79 -> fptosi -> muli -> %89 vs VECTOR +// store %extracted_12 -> fptosi -> muli -> %35). Rewriting the reference makes +// the VECTOR chain dead so retainNeededOpsInScope can drop it (and the +// math.floor/math.ceil feeding it), which is required for +// AnalyzeCubeControlFlowInputChain not to reject the module. +static void replaceVectorRefsInCubeScope(scope::ScopeOp cubeScope, + scope::ScopeOp vecScope) { + // Map (block_id, op_name) of a VECTOR-side math op to its transfer_id by + // walking the store value back to the tensor.extract tensor operand. + std::map, int64_t> vecMathToTid; + vecScope.walk([&](memref::StoreOp storeOp) { + auto tidAttr = + storeOp->getAttrOfType(CVPipeline::kTransferId); + if (!tidAttr) { + return; + } + Value storeVal = storeOp.getValue(); + if (auto ext = dyn_cast(storeVal.getDefiningOp())) { + Operation *tensorDef = ext.getTensor().getDefiningOp(); + if (tensorDef) { + auto blockIdOpt = CVPipeline::getOpBlockId(tensorDef); + if (blockIdOpt) { + vecMathToTid[{*blockIdOpt, + tensorDef->getName().getStringRef().str()}] = + tidAttr.getInt(); + } + } + } + }); + + // CUBE-side chains rooted at SSBuffer loads: (transfer, path) -> value. + std::map> cubePathsByTid; + cubeScope.walk([&](memref::LoadOp loadOp) { + auto tidAttr = + loadOp->getAttrOfType(CVPipeline::kTransferId); + if (!tidAttr) { + return; + } + llvm::DenseMap paths; + collectChainPaths(loadOp.getResult(), paths); + for (auto &pk : paths) { + cubePathsByTid[tidAttr.getInt()][pk.second] = pk.first; + } + }); + + // Boundary scalars in the CUBE scope derived from VECTOR math ops: value -> + // (transfer, path). The path is measured from the extracted scalar, so it + // aligns with the load-rooted chain path (both skip tensor.extract). + llvm::DenseMap> cubeVecBoundary; + cubeScope.walk([&](Operation *op) { + if (!isa(op->getDialect())) { + return; + } + auto blockIdOpt = CVPipeline::getOpBlockId(op); + if (!blockIdOpt) { + return; + } + std::string opName = op->getName().getStringRef().str(); + auto mIt = vecMathToTid.find({*blockIdOpt, opName}); + if (mIt == vecMathToTid.end()) { + return; + } + for (Value result : op->getResults()) { + for (Operation *user : result.getUsers()) { + auto ext = dyn_cast(user); + if (!ext) { + continue; + } + llvm::DenseMap paths; + collectChainPaths(ext.getResult(), paths); + for (auto &pk : paths) { + cubeVecBoundary[pk.first] = {mIt->second, pk.second}; + } + } + } + }); + + llvm::SmallVector> rewrites; + cubeScope.walk([&](Operation *op) { + for (mlir::OpOperand &operand : op->getOpOperands()) { + Value v = operand.get(); + auto bIt = cubeVecBoundary.find(v); + if (bIt == cubeVecBoundary.end()) { + continue; + } + int64_t tid = bIt->second.first; + const std::string &path = bIt->second.second; + auto tIt = cubePathsByTid.find(tid); + if (tIt == cubePathsByTid.end()) { + continue; + } + auto pIt = tIt->second.find(path); + if (pIt == tIt->second.end()) { + continue; + } + mlir::Value replacement = pIt->second; + if (replacement.getType() != v.getType()) { + continue; + } + rewrites.push_back({&operand, replacement}); + } + }); + for (auto &rw : rewrites) { + rw.first->set(rw.second); + } +} + +// Keep only the ops a scope actually needs, erasing the rest if trivially dead. +// +// After InterCoreTransferAndSync replaces tensor.extract results with SSBuffer +// loads, the extract op and its upstream VECTOR-only chain (math.floor/ceil, +// boundary computation) may be left behind in the CUBE scope. Each op in that +// chain references the next one, so a plain use_empty() check never triggers +// and the VECTOR-only ops keep leaking into the CUBE scope (later rejected by +// AnalyzeCubeControlFlowInputChain). +// +// Seeds the retained set with ops whose core_type matches the scope, then +// transitively retains every op feeding a retained op. Everything else that is +// trivially dead gets erased. +static void retainNeededOpsInScope(scope::ScopeOp scopeOp, StringRef scopeType) { + llvm::SmallVector ops; + scopeOp.walk([&](Operation *op) { + if (op->getNumRegions() == 0) { + ops.push_back(op); + } + }); + + llvm::DenseSet retained; + for (Operation *op : ops) { + if (matchesScope(op, scopeType)) { + retained.insert(op); + } + } + + bool changed = true; + while (changed) { + changed = false; + for (Operation *op : ops) { + if (retained.contains(op)) { + continue; + } + for (Value result : op->getResults()) { + for (Operation *user : result.getUsers()) { + if (retained.contains(user)) { + retained.insert(op); + changed = true; + break; + } + } + if (retained.contains(op)) { + break; + } + } + } + } + + // Iteratively erase use-free ops. This includes ops that matched the scope + // (retained as seeds): a tensor.extract that used to feed a cross-core scalar + // may have had all its uses rewritten to the SSBuffer load and is now dead; + // keeping it would keep the VECTOR-only math op it consumes alive in this + // scope. Ops with side effects (memref/llvm ops, etc.) are left alone by + // wouldOpBeTriviallyDead. + bool erased = true; + while (erased) { + erased = false; + llvm::SmallVector toErase; + scopeOp.walk([&](Operation *op) { + if (op->getNumRegions() > 0) { + return; + } + if (op->use_empty() && mlir::wouldOpBeTriviallyDead(op)) { + toErase.push_back(op); + } + }); + for (Operation *op : toErase) { + if (op->getBlock() && op->use_empty()) { + op->erase(); + erased = true; + } + } + } +} + static LogicalResult separateScopes(func::FuncOp funcOp) { debugDumpOperation("before SeparateCVScope on func", funcOp.getOperation()); @@ -974,6 +1207,13 @@ static LogicalResult separateScopes(func::FuncOp funcOp) { return failure(); } + // Rewrite leftover VECTOR-boundary references in the CUBE scope to their + // CUBE-side equivalents, then drop the now-dead chains so VECTOR-only ops do + // not leak into the CUBE scope. + replaceVectorRefsInCubeScope(cubeScope, vecScope); + retainNeededOpsInScope(cubeScope, "CUBE"); + retainNeededOpsInScope(vecScope, "VECTOR"); + cleanupSsbufferAttrs(funcOp); debugDumpOperation("after SeparateCVScope on func", funcOp.getOperation()); @@ -1015,6 +1255,58 @@ void mlir::triton::SeparateCVScopePass::runOnOperation() { UnitAttr::get(scopeOp->getContext())); }); + // Eliminate redundant store→load pairs in VECTOR scopes. + // InterCoreTransferAndSync inserts llvm.store (send to SSBuffer) followed by + // llvm.load (receive on the consumer side). After scope separation the + // VECTOR scope contains both: the store sends the value, and the load re-reads + // it for use as a for-loop bound. The load is redundant — the stored value + // is already live — so replace loaded values with the stored value. + module.walk([](scope::ScopeOp scopeOp) { + auto coreTypeAttr = + scopeOp->getAttrOfType(hivm::TCoreTypeAttr::name); + if (!coreTypeAttr || coreTypeAttr.getTcoretype() != hivm::TCoreType::VECTOR) { + return; + } + + llvm::DenseMap storedValues; + scopeOp.walk([&](memref::StoreOp storeOp) { + auto transferIdAttr = + storeOp->getAttrOfType(CVPipeline::kTransferId); + if (!transferIdAttr) { + return; + } + int64_t tid = transferIdAttr.getInt(); + storedValues[tid] = storeOp.getValue(); + }); + + if (storedValues.empty()) { + return; + } + + llvm::SmallVector deadLoads; + scopeOp.walk([&](memref::LoadOp loadOp) { + auto transferIdAttr = + loadOp->getAttrOfType(CVPipeline::kTransferId); + if (!transferIdAttr) { + return; + } + int64_t tid = transferIdAttr.getInt(); + auto it = storedValues.find(tid); + if (it == storedValues.end()) { + return; + } + mlir::Value storeVal = it->second; + if (storeVal == loadOp.getResult()) { + return; + } + loadOp.replaceAllUsesWith(storeVal); + deadLoads.push_back(loadOp); + }); + for (memref::LoadOp loadOp : deadLoads) { + loadOp->erase(); + } + }); + debugDumpOperation("after SeparateCVScopePass", module.getOperation()); } From 40e0cbe6606ec292388c66358bfa04678b265946 Mon Sep 17 00:00:00 2001 From: 1801ZDL <1241272204@qq.com> Date: Wed, 12 Aug 2026 16:42:39 +0800 Subject: [PATCH 2/9] fix: classify scalar-dep extracts as VECTOR; split extract analysis Review feedback: 1. OpClassifier: an extract whose source tensor traces to a VECTOR-only producer (e.g. math.floor/ceil) is itself a VECTOR computation -- the CUBE side consumes its scalar via the SSBuffer dependency channel and must not recompute it. Skip marking such extracts CUBE in both propagateCubeUpstream and propagateCubeUpstreamForOp (new hasVectorOnlyProducer helper). 2. DataDependencyAnalysis: split the tensor.extract walk out of analyzeScalarVToCDependencies into analyzeScalarExtractDependencies, since the 3-walk function was ~300 lines. Pure refactor, no behavior change. Co-Authored-By: DeLong code --- .../SplitDataflow/DataDependencyAnalysis.h | 4 ++ .../PlanComputeBlock/OpClassifier.cpp | 43 ++++++++++++ .../SplitDataflow/DataDependencyAnalysis.cpp | 65 +++++++++++-------- 3 files changed, 86 insertions(+), 26 deletions(-) diff --git a/third_party/ascend/include/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.h b/third_party/ascend/include/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.h index b2392219df..f6c21d8dbe 100644 --- a/third_party/ascend/include/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.h +++ b/third_party/ascend/include/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.h @@ -145,6 +145,10 @@ class DataDependencyAnalysisPass void analyzeExternalInputs(DataDependencyInfo &info); void analyzeExternalOutputs(DataDependencyInfo &info); void analyzeScalarVToCDependencies(DataDependencyInfo &info); + void analyzeScalarExtractDependencies( + DataDependencyInfo &info, + llvm::DenseSet &handledScalarValues, + llvm::DenseSet &handledForOps); void analyzeMemoryEffect(DataDependencyInfo &info); std::pair findCommonLevelBlockIds(DataDependencyInfo &info, diff --git a/third_party/ascend/lib/DynamicCVPipeline/PlanComputeBlock/OpClassifier.cpp b/third_party/ascend/lib/DynamicCVPipeline/PlanComputeBlock/OpClassifier.cpp index 34e1e18ecb..e402663259 100644 --- a/third_party/ascend/lib/DynamicCVPipeline/PlanComputeBlock/OpClassifier.cpp +++ b/third_party/ascend/lib/DynamicCVPipeline/PlanComputeBlock/OpClassifier.cpp @@ -717,6 +717,30 @@ static bool isTensorArithOrMathOp(Operation *op) { return false; } +// True if `value`'s defining chain reaches a VECTOR-only op. A tensor.extract +// whose source tensor is produced by such an op (e.g. math.floor/ceil) is +// itself a VECTOR computation: the CUBE side consumes its scalar via the +// SSBuffer dependency channel and must not recompute the extract, so it should +// not be marked CUBE. +static bool hasVectorOnlyProducer(Value value) { + llvm::SmallVector worklist{value}; + llvm::DenseSet visited; + while (!worklist.empty()) { + Value cur = worklist.pop_back_val(); + Operation *defOp = cur.getDefiningOp(); + if (!defOp || !visited.insert(defOp).second) { + continue; + } + if (CVPipeline::isVectorOnlyOp(defOp)) { + return true; + } + for (Value operand : defOp->getOperands()) { + worklist.push_back(operand); + } + } + return false; +} + // Propagate CUBE core type upstream int OpClassifierPass::propagateCubeUpstream() { LLVM_DEBUG(DBGS() << "--- Step 2: CUBE upstream BFS --->\n"); @@ -751,6 +775,18 @@ int OpClassifierPass::propagateCubeUpstream() { continue; } + // An extract of a VECTOR-only-produced tensor is itself a VECTOR + // computation; the CUBE side consumes its scalar via SSBuffer and must + // not be marked CUBE. + if (auto extOp = dyn_cast(def)) { + if (hasVectorOnlyProducer(extOp.getTensor())) { + cubeVisited.insert(def); + LLVM_DEBUG(DBGS() << "skip " << def->getName().getStringRef() + << ": extract of vector-only producer\n"); + continue; + } + } + // Skip operations inside linalg block (internal values) // But don't skip the linalg op itself if (isInsideNestedLinalgRegion(def)) { @@ -938,6 +974,13 @@ void OpClassifierPass::propagateCubeUpstreamForOp(Operation *startOp) { if (isTensorArithOrMathOp(upstreamOp)) continue; + // Extract of a VECTOR-only-produced tensor is a VECTOR computation; the + // CUBE side consumes its scalar via SSBuffer, not by recomputing it. + if (auto extOp = dyn_cast(upstreamOp)) { + if (hasVectorOnlyProducer(extOp.getTensor())) + continue; + } + cubeVisited.insert(upstreamOp); LLVM_DEBUG(DBGS() << "\t\tcube upstream: " << upstreamOp->getName().getStringRef() << "\n"); diff --git a/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.cpp b/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.cpp index 3f55880e1c..ae1587dd32 100644 --- a/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.cpp +++ b/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.cpp @@ -918,35 +918,16 @@ static bool ifOpHasCubeOps( return hasCube; } -// Analyze scalar V->C dependencies from control flow ops. -// Detects when scf.for loop bounds or scf.if conditions are scalar values whose -// defining chain traces back to vector-only ops (e.g. math.floor/math.ceil on -// tensors, linalg.reduce). These scalars must be transferred from VECTOR to CUBE. -// -// To avoid redundant transfers, scalars already recorded as dependencies are -// tracked in a stop-set: downstream scalars whose defining chain only reaches -// vector-only ops through an already-handled scalar are NOT re-recorded. -void DataDependencyAnalysisPass::analyzeScalarVToCDependencies( - DataDependencyInfo &info) { +// Detect scalars extracted from VECTOR-produced tensors and consumed by CUBE +// blocks. The extract result is the natural scalar dependency boundary: +// transferring it via SSBuffer avoids the need for 1-D tensor CopyOps. +void DataDependencyAnalysisPass::analyzeScalarExtractDependencies( + DataDependencyInfo &info, + llvm::DenseSet &handledScalarValues, + llvm::DenseSet &handledForOps) { auto &blockInfoMap = info.getBlockInfoMap(); auto &v2cDependencies = info.getV2CDependencies(); - // Scalars already recorded as V->C deps — downstream values depending on - // these don't need separate transfer since the upstream value will be - // available on the CUBE side after SSBuffer transfer. - llvm::DenseSet handledScalarValues; - - // For-loops whose bounds have been handled — ifOp conditions inside these - // loops will be computable on CUBE via the induction variable after the - // bounds are transferred, so they don't need separate scalar deps. - llvm::DenseSet handledForOps; - - LOG_DEBUG("Analyzing scalar V->C dependencies from control flow ops...\n"); - - // ---- tensor.extract ops ---- - // Detect scalars extracted from VECTOR-produced tensors and consumed by CUBE - // blocks. The extract result is the natural scalar dependency boundary: - // transferring it via SSBuffer avoids the need for 1-D tensor CopyOps. module.walk([&](tensor::ExtractOp extractOp) { mlir::Value sourceTensor = extractOp.getTensor(); mlir::Operation *tensorDefOp = sourceTensor.getDefiningOp(); @@ -1063,6 +1044,38 @@ void DataDependencyAnalysisPass::analyzeScalarVToCDependencies( } } }); +} + +// Analyze scalar V->C dependencies from control flow ops. +// Detects when scf.for loop bounds or scf.if conditions are scalar values whose +// defining chain traces back to vector-only ops (e.g. math.floor/math.ceil on +// tensors, linalg.reduce). These scalars must be transferred from VECTOR to CUBE. +// +// To avoid redundant transfers, scalars already recorded as dependencies are +// tracked in a stop-set: downstream scalars whose defining chain only reaches +// vector-only ops through an already-handled scalar are NOT re-recorded. +void DataDependencyAnalysisPass::analyzeScalarVToCDependencies( + DataDependencyInfo &info) { + auto &blockInfoMap = info.getBlockInfoMap(); + auto &v2cDependencies = info.getV2CDependencies(); + + // Scalars already recorded as V->C deps — downstream values depending on + // these don't need separate transfer since the upstream value will be + // available on the CUBE side after SSBuffer transfer. + llvm::DenseSet handledScalarValues; + + // For-loops whose bounds have been handled — ifOp conditions inside these + // loops will be computable on CUBE via the induction variable after the + // bounds are transferred, so they don't need separate scalar deps. + llvm::DenseSet handledForOps; + + LOG_DEBUG("Analyzing scalar V->C dependencies from control flow ops...\n"); + + // ---- tensor.extract ops ---- + // Detect scalars extracted from VECTOR-produced tensors and consumed by CUBE + // blocks. The extract result is the natural scalar dependency boundary: + // transferring it via SSBuffer avoids the need for 1-D tensor CopyOps. + analyzeScalarExtractDependencies(info, handledScalarValues, handledForOps); // ---- scf.for loop bounds ---- module.walk([&](scf::ForOp forOp) { From e4470c8d9c365d13997d24691e1d9a9b7b8f976a Mon Sep 17 00:00:00 2001 From: 1801ZDL <1241272204@qq.com> Date: Wed, 12 Aug 2026 17:11:40 +0800 Subject: [PATCH 3/9] fix: drop extract-in-CUBE-block hack; dedup vs external-input deps Follow-up to classifying scalar-dep extracts as VECTOR: - With the extract now VECTOR, requiring it to live in a CUBE block would reject the very extracts the scalar V->C feature targets. Remove that block-level isCube gate and rely on the downstream CUBE-consumer walk. - analyzeExternalInputs (which runs first) may already have recorded a V->C dep for the extract crossing into a CUBE block. Skip the duplicate in analyzeScalarExtractDependencies, otherwise the same scalar gets a second SSBuffer store/sync (deadlock risk). Verified end-to-end from the pre-plan input: op-classifier now classifies the floor/ceil extracts as VECTOR, the scalar transfer is inserted with one store per scalar, and the CUBE scope is free of math ops. Co-Authored-By: DeLong code --- .../SplitDataflow/DataDependencyAnalysis.cpp | 30 ++++++++++--------- 1 file changed, 16 insertions(+), 14 deletions(-) diff --git a/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.cpp b/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.cpp index ae1587dd32..321b6b0f05 100644 --- a/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.cpp +++ b/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.cpp @@ -949,20 +949,14 @@ void DataDependencyAnalysisPass::analyzeScalarExtractDependencies( return; } - // Only handle extracts located in a CUBE block. The VECTOR-side extract - // (e.g. block 21) is the original computation; the CUBE-side extract - // (e.g. block 15) is the copy whose scalar result the CUBE side actually - // consumes. Transferring the VECTOR-side extract creates a CUBE load with - // no users (redundant transfer). + // The extract itself is now classified VECTOR by OpClassifier (its source + // tensor traces to a VECTOR-only producer), so it is no longer restricted + // to CUBE blocks. Whether the CUBE side actually needs the scalar is + // decided by the downstream CUBE-consumer walk below. auto extractBlockIdOpt = CVPipeline::getOpBlockId(extractOp.getOperation()); if (!extractBlockIdOpt) { return; } - auto extractBlockIt = blockInfoMap.find(*extractBlockIdOpt); - if (extractBlockIt == blockInfoMap.end() || - !extractBlockIt->second.isCube) { - return; - } mlir::Value scalarResult = extractOp.getResult(); if (!isa(scalarResult.getType())) { @@ -1020,10 +1014,18 @@ void DataDependencyAnalysisPass::analyzeScalarExtractDependencies( } } if (hasCubConsumer) { - // Create a single dependency for this extract result. All consumers - // share one SSBuffer store; one dep avoids duplicate stores (and the - // duplicate sync that can deadlock the cores). - if (!handledConsumers.empty()) { + // analyzeExternalInputs runs before this walk and may already have + // recorded a V->C dependency for this extract (it crosses into a CUBE + // block as an external input). Skip the duplicate: a second dep would + // produce a second SSBuffer store/sync for the same value and can + // deadlock the cores. + bool alreadyDep = llvm::any_of(v2cDependencies, [&](const DependencyInfo &d) { + return d.value == scalarResult; + }); + if (!alreadyDep && !handledConsumers.empty()) { + // Create a single dependency for this extract result. All consumers + // share one SSBuffer store; one dep avoids duplicate stores (and the + // duplicate sync that can deadlock the cores). int consumerId = *handledConsumers.begin(); collectDepInfo(scalarResult, DependencyType::VectorToCube, v2cDependencies, producerId, consumerId, info); From 874d73bfcbdf06f74d2750efe89859d940bf8c9d Mon Sep 17 00:00:00 2001 From: 1801ZDL <1241272204@qq.com> Date: Wed, 12 Aug 2026 17:22:37 +0800 Subject: [PATCH 4/9] fix: disable CUBE clone-chain cleanup pending E2E verification With OpClassifier now marking scalar-dependency extracts (and their math chains) as VECTOR, the CUBE scope no longer contains these chains, so the replaceVectorRefsInCubeScope/retainNeededOpsInScope rewrite+cleanup is expected to be dead code. Comment the calls out for an end-to-end test; the definitions are removed once E2E passes. Co-Authored-By: DeLong code --- .../SplitDataflow/SeparateCVScope.cpp | 11 ++++++++--- 1 file changed, 8 insertions(+), 3 deletions(-) diff --git a/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/SeparateCVScope.cpp b/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/SeparateCVScope.cpp index cd61c417cb..8606e8ca35 100644 --- a/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/SeparateCVScope.cpp +++ b/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/SeparateCVScope.cpp @@ -1210,9 +1210,14 @@ static LogicalResult separateScopes(func::FuncOp funcOp) { // Rewrite leftover VECTOR-boundary references in the CUBE scope to their // CUBE-side equivalents, then drop the now-dead chains so VECTOR-only ops do // not leak into the CUBE scope. - replaceVectorRefsInCubeScope(cubeScope, vecScope); - retainNeededOpsInScope(cubeScope, "CUBE"); - retainNeededOpsInScope(vecScope, "VECTOR"); + // + // TEMP: disabled pending end-to-end verification. With OpClassifier now + // marking scalar-dependency extracts (and their math chains) as VECTOR, the + // CUBE scope no longer contains these chains, so this rewrite/cleanup is + // expected to be dead code. To be removed if the E2E test passes. + // replaceVectorRefsInCubeScope(cubeScope, vecScope); + // retainNeededOpsInScope(cubeScope, "CUBE"); + // retainNeededOpsInScope(vecScope, "VECTOR"); cleanupSsbufferAttrs(funcOp); From 844afc3d71c8f3fed016cf13ecc47a6a4bfd44b6 Mon Sep 17 00:00:00 2001 From: 1801ZDL <1241272204@qq.com> Date: Wed, 12 Aug 2026 17:36:25 +0800 Subject: [PATCH 5/9] refactor: remove now-dead CUBE clone-chain cleanup (mechanism B) E2E verified that with OpClassifier marking scalar-dependency extracts (and their math chains) as VECTOR, the CUBE scope no longer contains the cloned VECTOR chains, so the rewrite+cleanup in SeparateCVScope is dead code. Remove collectChainPaths, replaceVectorRefsInCubeScope and retainNeededOpsInScope, their call sites, and the now-unused includes (LLVM/Tensor) plus duplicate includes. Co-Authored-By: DeLong code --- .../SplitDataflow/SeparateCVScope.cpp | 248 +----------------- 1 file changed, 1 insertion(+), 247 deletions(-) diff --git a/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/SeparateCVScope.cpp b/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/SeparateCVScope.cpp index 8606e8ca35..d5a366160a 100644 --- a/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/SeparateCVScope.cpp +++ b/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/SeparateCVScope.cpp @@ -29,14 +29,13 @@ #include "llvm/Support/Debug.h" #include "ascend/include/DynamicCVPipeline/Common/Utils.h" +#include "ascend/include/DynamicCVPipeline/SplitDataflow/SeparateCVScope.h" #include "bishengir/Dialect/HIVM/IR/HIVM.h" #include "bishengir/Dialect/Scope/IR/Scope.h" #include "mlir/Dialect/Arith/IR/Arith.h" #include "mlir/Dialect/Func/IR/FuncOps.h" -#include "mlir/Dialect/LLVMIR/LLVMDialect.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" #include "mlir/Dialect/SCF/IR/SCF.h" -#include "mlir/Dialect/Tensor/IR/Tensor.h" #include "mlir/IR/Builders.h" #include "mlir/IR/BuiltinAttributes.h" #include "mlir/IR/IRMapping.h" @@ -44,12 +43,6 @@ #include "mlir/Pass/Pass.h" #include "mlir/Transforms/RegionUtils.h" -#include "ascend/include/DynamicCVPipeline/Common/Utils.h" -#include "ascend/include/DynamicCVPipeline/SplitDataflow/SeparateCVScope.h" - -#include "bishengir/Dialect/HIVM/IR/HIVM.h" -#include "bishengir/Dialect/Scope/IR/Scope.h" - using namespace mlir; using namespace mlir::triton; @@ -956,233 +949,6 @@ static void cleanupSsbufferAttrs(Operation *rootOp) { rootOp->walk([](Operation *op) { removeSsbufferAttrs(op); }); } -// Record, for every value produced along the forward use-chain rooted at -// `root`, the sequence of op names that led to it. This lets us pair up a -// value on the VECTOR side (from a SSBuffer store) with the structurally -// identical value on the CUBE side (from the matching SSBuffer load). -static void collectChainPaths( - Value root, llvm::DenseMap &paths) { - llvm::SmallVector worklist; - worklist.push_back(root); - paths[root] = ""; - while (!worklist.empty()) { - Value cur = worklist.pop_back_val(); - std::string curPath = paths[cur]; - Region *rootRegion = root.getParentRegion(); - for (Operation *user : cur.getUsers()) { - if (user->getNumResults() == 0) { - continue; - } - // Only follow single-result pure-compute ops; stop at anything with - // regions or side effects. - if (user->getNumRegions() > 0 || - !mlir::wouldOpBeTriviallyDead(user)) { - continue; - } - // Do not cross into nested regions (e.g. a scf.for body): a use inside a - // loop is not a candidate for the boundary-mapping used here, and - // substituting a value defined there would break dominance. - if (user->getBlock()->getParent() != rootRegion) { - continue; - } - Value res = user->getResult(0); - if (paths.contains(res)) { - continue; - } - paths[res] = curPath + user->getName().getStringRef().str() + ";"; - worklist.push_back(res); - } - } -} - -// Rewrite uses in the CUBE scope that reference values computed on the VECTOR -// side (boundary scalars derived from math.floor/math.ceil via tensor.extract). -// -// Such references are left-over copies of the VECTOR boundary chain that -// SeparateCVScope cloned into the CUBE scope's mixed for-loop. The CUBE scope -// already has an equivalent chain rooted at the SSBuffer load for the same -// transfer_id (e.g. load %79 -> fptosi -> muli -> %89 vs VECTOR -// store %extracted_12 -> fptosi -> muli -> %35). Rewriting the reference makes -// the VECTOR chain dead so retainNeededOpsInScope can drop it (and the -// math.floor/math.ceil feeding it), which is required for -// AnalyzeCubeControlFlowInputChain not to reject the module. -static void replaceVectorRefsInCubeScope(scope::ScopeOp cubeScope, - scope::ScopeOp vecScope) { - // Map (block_id, op_name) of a VECTOR-side math op to its transfer_id by - // walking the store value back to the tensor.extract tensor operand. - std::map, int64_t> vecMathToTid; - vecScope.walk([&](memref::StoreOp storeOp) { - auto tidAttr = - storeOp->getAttrOfType(CVPipeline::kTransferId); - if (!tidAttr) { - return; - } - Value storeVal = storeOp.getValue(); - if (auto ext = dyn_cast(storeVal.getDefiningOp())) { - Operation *tensorDef = ext.getTensor().getDefiningOp(); - if (tensorDef) { - auto blockIdOpt = CVPipeline::getOpBlockId(tensorDef); - if (blockIdOpt) { - vecMathToTid[{*blockIdOpt, - tensorDef->getName().getStringRef().str()}] = - tidAttr.getInt(); - } - } - } - }); - - // CUBE-side chains rooted at SSBuffer loads: (transfer, path) -> value. - std::map> cubePathsByTid; - cubeScope.walk([&](memref::LoadOp loadOp) { - auto tidAttr = - loadOp->getAttrOfType(CVPipeline::kTransferId); - if (!tidAttr) { - return; - } - llvm::DenseMap paths; - collectChainPaths(loadOp.getResult(), paths); - for (auto &pk : paths) { - cubePathsByTid[tidAttr.getInt()][pk.second] = pk.first; - } - }); - - // Boundary scalars in the CUBE scope derived from VECTOR math ops: value -> - // (transfer, path). The path is measured from the extracted scalar, so it - // aligns with the load-rooted chain path (both skip tensor.extract). - llvm::DenseMap> cubeVecBoundary; - cubeScope.walk([&](Operation *op) { - if (!isa(op->getDialect())) { - return; - } - auto blockIdOpt = CVPipeline::getOpBlockId(op); - if (!blockIdOpt) { - return; - } - std::string opName = op->getName().getStringRef().str(); - auto mIt = vecMathToTid.find({*blockIdOpt, opName}); - if (mIt == vecMathToTid.end()) { - return; - } - for (Value result : op->getResults()) { - for (Operation *user : result.getUsers()) { - auto ext = dyn_cast(user); - if (!ext) { - continue; - } - llvm::DenseMap paths; - collectChainPaths(ext.getResult(), paths); - for (auto &pk : paths) { - cubeVecBoundary[pk.first] = {mIt->second, pk.second}; - } - } - } - }); - - llvm::SmallVector> rewrites; - cubeScope.walk([&](Operation *op) { - for (mlir::OpOperand &operand : op->getOpOperands()) { - Value v = operand.get(); - auto bIt = cubeVecBoundary.find(v); - if (bIt == cubeVecBoundary.end()) { - continue; - } - int64_t tid = bIt->second.first; - const std::string &path = bIt->second.second; - auto tIt = cubePathsByTid.find(tid); - if (tIt == cubePathsByTid.end()) { - continue; - } - auto pIt = tIt->second.find(path); - if (pIt == tIt->second.end()) { - continue; - } - mlir::Value replacement = pIt->second; - if (replacement.getType() != v.getType()) { - continue; - } - rewrites.push_back({&operand, replacement}); - } - }); - for (auto &rw : rewrites) { - rw.first->set(rw.second); - } -} - -// Keep only the ops a scope actually needs, erasing the rest if trivially dead. -// -// After InterCoreTransferAndSync replaces tensor.extract results with SSBuffer -// loads, the extract op and its upstream VECTOR-only chain (math.floor/ceil, -// boundary computation) may be left behind in the CUBE scope. Each op in that -// chain references the next one, so a plain use_empty() check never triggers -// and the VECTOR-only ops keep leaking into the CUBE scope (later rejected by -// AnalyzeCubeControlFlowInputChain). -// -// Seeds the retained set with ops whose core_type matches the scope, then -// transitively retains every op feeding a retained op. Everything else that is -// trivially dead gets erased. -static void retainNeededOpsInScope(scope::ScopeOp scopeOp, StringRef scopeType) { - llvm::SmallVector ops; - scopeOp.walk([&](Operation *op) { - if (op->getNumRegions() == 0) { - ops.push_back(op); - } - }); - - llvm::DenseSet retained; - for (Operation *op : ops) { - if (matchesScope(op, scopeType)) { - retained.insert(op); - } - } - - bool changed = true; - while (changed) { - changed = false; - for (Operation *op : ops) { - if (retained.contains(op)) { - continue; - } - for (Value result : op->getResults()) { - for (Operation *user : result.getUsers()) { - if (retained.contains(user)) { - retained.insert(op); - changed = true; - break; - } - } - if (retained.contains(op)) { - break; - } - } - } - } - - // Iteratively erase use-free ops. This includes ops that matched the scope - // (retained as seeds): a tensor.extract that used to feed a cross-core scalar - // may have had all its uses rewritten to the SSBuffer load and is now dead; - // keeping it would keep the VECTOR-only math op it consumes alive in this - // scope. Ops with side effects (memref/llvm ops, etc.) are left alone by - // wouldOpBeTriviallyDead. - bool erased = true; - while (erased) { - erased = false; - llvm::SmallVector toErase; - scopeOp.walk([&](Operation *op) { - if (op->getNumRegions() > 0) { - return; - } - if (op->use_empty() && mlir::wouldOpBeTriviallyDead(op)) { - toErase.push_back(op); - } - }); - for (Operation *op : toErase) { - if (op->getBlock() && op->use_empty()) { - op->erase(); - erased = true; - } - } - } -} static LogicalResult separateScopes(func::FuncOp funcOp) { debugDumpOperation("before SeparateCVScope on func", funcOp.getOperation()); @@ -1207,18 +973,6 @@ static LogicalResult separateScopes(func::FuncOp funcOp) { return failure(); } - // Rewrite leftover VECTOR-boundary references in the CUBE scope to their - // CUBE-side equivalents, then drop the now-dead chains so VECTOR-only ops do - // not leak into the CUBE scope. - // - // TEMP: disabled pending end-to-end verification. With OpClassifier now - // marking scalar-dependency extracts (and their math chains) as VECTOR, the - // CUBE scope no longer contains these chains, so this rewrite/cleanup is - // expected to be dead code. To be removed if the E2E test passes. - // replaceVectorRefsInCubeScope(cubeScope, vecScope); - // retainNeededOpsInScope(cubeScope, "CUBE"); - // retainNeededOpsInScope(vecScope, "VECTOR"); - cleanupSsbufferAttrs(funcOp); debugDumpOperation("after SeparateCVScope on func", funcOp.getOperation()); From a6fcf381ed2aa0f5958a95e53316a8b288ecf12d Mon Sep 17 00:00:00 2001 From: 1801ZDL <1241272204@qq.com> Date: Wed, 12 Aug 2026 17:47:08 +0800 Subject: [PATCH 6/9] style: condense verbose comments in scalar V->C feature Shorten the added comments (no examples/walkthroughs), and fix a stale llvm.store/load reference in the VECTOR redundant-load cleanup comment. Co-Authored-By: DeLong code --- .../Common/SSBufferManager.cpp | 4 +- .../PlanComputeBlock/OpClassifier.cpp | 23 +++---- .../SplitDataflow/DataDependencyAnalysis.cpp | 68 ++++++------------- .../InterCoreTransferAndSync.cpp | 18 ++--- .../SplitDataflow/SeparateCVScope.cpp | 8 +-- 5 files changed, 39 insertions(+), 82 deletions(-) diff --git a/third_party/ascend/lib/DynamicCVPipeline/Common/SSBufferManager.cpp b/third_party/ascend/lib/DynamicCVPipeline/Common/SSBufferManager.cpp index 4a4761f8a5..a5c383ccb6 100644 --- a/third_party/ascend/lib/DynamicCVPipeline/Common/SSBufferManager.cpp +++ b/third_party/ascend/lib/DynamicCVPipeline/Common/SSBufferManager.cpp @@ -103,9 +103,7 @@ SSBufferManager::writeToSSBuffer(Value value, OpBuilder &builder, int64_t addrValue = addrResult.value(); Location loc = builder.getUnknownLoc(); - // The memref element type must match the stored value's type: with LLVM - // pointer ops the element type was untyped, but memref.store verifies that - // the value type matches the memref element type. + // memref.store requires value type == memref element type. auto [constOp, pointerCastOp] = getSsbufConstAndPointerCast(builder, loc, addrValue, value.getType()); createdOps.push_back(constOp); diff --git a/third_party/ascend/lib/DynamicCVPipeline/PlanComputeBlock/OpClassifier.cpp b/third_party/ascend/lib/DynamicCVPipeline/PlanComputeBlock/OpClassifier.cpp index e402663259..dadaa248bc 100644 --- a/third_party/ascend/lib/DynamicCVPipeline/PlanComputeBlock/OpClassifier.cpp +++ b/third_party/ascend/lib/DynamicCVPipeline/PlanComputeBlock/OpClassifier.cpp @@ -703,8 +703,7 @@ void OpClassifierPass::getUpstreamOpsWithMemoryDeps( } } -// An arith/math op with a tensor result is VECTOR-only and must not be marked -// CUBE; scalar arith/math (index/i32 computation) may still be marked CUBE. +// arith/math op with a tensor result is VECTOR-only (not CUBE). static bool isTensorArithOrMathOp(Operation *op) { if (!isa(op->getDialect())) { return false; @@ -717,11 +716,9 @@ static bool isTensorArithOrMathOp(Operation *op) { return false; } -// True if `value`'s defining chain reaches a VECTOR-only op. A tensor.extract -// whose source tensor is produced by such an op (e.g. math.floor/ceil) is -// itself a VECTOR computation: the CUBE side consumes its scalar via the -// SSBuffer dependency channel and must not recompute the extract, so it should -// not be marked CUBE. +// True if `value`'s defining chain reaches a VECTOR-only op. An extract of +// such a tensor is itself VECTOR (CUBE gets the scalar via SSBuffer), so it +// must not be marked CUBE. static bool hasVectorOnlyProducer(Value value) { llvm::SmallVector worklist{value}; llvm::DenseSet visited; @@ -775,9 +772,7 @@ int OpClassifierPass::propagateCubeUpstream() { continue; } - // An extract of a VECTOR-only-produced tensor is itself a VECTOR - // computation; the CUBE side consumes its scalar via SSBuffer and must - // not be marked CUBE. + // An extract of a VECTOR-only tensor is itself VECTOR. if (auto extOp = dyn_cast(def)) { if (hasVectorOnlyProducer(extOp.getTensor())) { cubeVisited.insert(def); @@ -968,14 +963,12 @@ void OpClassifierPass::propagateCubeUpstreamForOp(Operation *startOp) { continue; if (isa(upstreamOp)) continue; - // Align with the main propagateCubeUpstream: only skip arith/math ops - // with tensor results (they are VECTOR-only); scalar arith/math ops - // (index/i32 computation) may still be marked CUBE. + // Align with propagateCubeUpstream: skip arith/math with tensor results + // (scalar arith/math may still be marked CUBE). if (isTensorArithOrMathOp(upstreamOp)) continue; - // Extract of a VECTOR-only-produced tensor is a VECTOR computation; the - // CUBE side consumes its scalar via SSBuffer, not by recomputing it. + // Extract of a VECTOR-only tensor is itself VECTOR. if (auto extOp = dyn_cast(upstreamOp)) { if (hasVectorOnlyProducer(extOp.getTensor())) continue; diff --git a/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.cpp b/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.cpp index 321b6b0f05..6d2ac3a825 100644 --- a/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.cpp +++ b/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.cpp @@ -918,9 +918,7 @@ static bool ifOpHasCubeOps( return hasCube; } -// Detect scalars extracted from VECTOR-produced tensors and consumed by CUBE -// blocks. The extract result is the natural scalar dependency boundary: -// transferring it via SSBuffer avoids the need for 1-D tensor CopyOps. +// Detect scalars extracted from VECTOR-only tensors and consumed by CUBE. void DataDependencyAnalysisPass::analyzeScalarExtractDependencies( DataDependencyInfo &info, llvm::DenseSet &handledScalarValues, @@ -949,10 +947,8 @@ void DataDependencyAnalysisPass::analyzeScalarExtractDependencies( return; } - // The extract itself is now classified VECTOR by OpClassifier (its source - // tensor traces to a VECTOR-only producer), so it is no longer restricted - // to CUBE blocks. Whether the CUBE side actually needs the scalar is - // decided by the downstream CUBE-consumer walk below. + // Whether the CUBE side needs the scalar is decided by the downstream + // CUBE-consumer walk below (the extract itself is VECTOR-classified). auto extractBlockIdOpt = CVPipeline::getOpBlockId(extractOp.getOperation()); if (!extractBlockIdOpt) { return; @@ -963,9 +959,8 @@ void DataDependencyAnalysisPass::analyzeScalarExtractDependencies( return; } - // The scalar must be consumed (directly or through a downstream pure - // scalar chain, e.g. fptosi/muli/subi producing loop bounds) in at least - // one CUBE block or in a for/if with CUBE content. + // The scalar must be consumed (directly or via a pure scalar chain) by at + // least one CUBE block or a for/if with CUBE content. int producerId = *extractBlockIdOpt; llvm::DenseSet handledConsumers; @@ -1014,18 +1009,13 @@ void DataDependencyAnalysisPass::analyzeScalarExtractDependencies( } } if (hasCubConsumer) { - // analyzeExternalInputs runs before this walk and may already have - // recorded a V->C dependency for this extract (it crosses into a CUBE - // block as an external input). Skip the duplicate: a second dep would - // produce a second SSBuffer store/sync for the same value and can - // deadlock the cores. + // analyzeExternalInputs may already have a dep for this extract; skip + // the duplicate (a second store/sync could deadlock the cores). bool alreadyDep = llvm::any_of(v2cDependencies, [&](const DependencyInfo &d) { return d.value == scalarResult; }); if (!alreadyDep && !handledConsumers.empty()) { - // Create a single dependency for this extract result. All consumers - // share one SSBuffer store; one dep avoids duplicate stores (and the - // duplicate sync that can deadlock the cores). + // One dep per extract; all consumers share a single SSBuffer store. int consumerId = *handledConsumers.begin(); collectDepInfo(scalarResult, DependencyType::VectorToCube, v2cDependencies, producerId, consumerId, info); @@ -1033,9 +1023,8 @@ void DataDependencyAnalysisPass::analyzeScalarExtractDependencies( LOG_DEBUG("Found scalar V->C dependency from tensor.extract: " << scalarResult << "\n"); handledScalarValues.insert(scalarResult); - // Mark any enclosing for-loop as handled: its bounds and any ifOp - // conditions inside derive from this transferred scalar, so they don't - // need separate transfers. + // Enclosing loop's bounds/conditions derive from this scalar: no further + // transfers needed for them. mlir::Operation *parent = extractOp->getParentOp(); while (parent) { if (auto parentForOp = dyn_cast(parent)) { @@ -1048,27 +1037,19 @@ void DataDependencyAnalysisPass::analyzeScalarExtractDependencies( }); } -// Analyze scalar V->C dependencies from control flow ops. -// Detects when scf.for loop bounds or scf.if conditions are scalar values whose -// defining chain traces back to vector-only ops (e.g. math.floor/math.ceil on -// tensors, linalg.reduce). These scalars must be transferred from VECTOR to CUBE. -// -// To avoid redundant transfers, scalars already recorded as dependencies are -// tracked in a stop-set: downstream scalars whose defining chain only reaches -// vector-only ops through an already-handled scalar are NOT re-recorded. +// Detect scalar V->C deps from extracts, for-loop bounds and if conditions +// whose defining chain reaches a vector-only op. A stop-set suppresses +// redundant transfers for scalars derived from already-handled ones. void DataDependencyAnalysisPass::analyzeScalarVToCDependencies( DataDependencyInfo &info) { auto &blockInfoMap = info.getBlockInfoMap(); auto &v2cDependencies = info.getV2CDependencies(); - // Scalars already recorded as V->C deps — downstream values depending on - // these don't need separate transfer since the upstream value will be - // available on the CUBE side after SSBuffer transfer. + // Scalars already transferred; downstream derivatives need no new dep. llvm::DenseSet handledScalarValues; - // For-loops whose bounds have been handled — ifOp conditions inside these - // loops will be computable on CUBE via the induction variable after the - // bounds are transferred, so they don't need separate scalar deps. + // For-loops whose bounds were transferred: inner ifOp conditions are then + // computable via the induction variable, so no separate deps are needed. llvm::DenseSet handledForOps; LOG_DEBUG("Analyzing scalar V->C dependencies from control flow ops...\n"); @@ -1108,11 +1089,8 @@ void DataDependencyAnalysisPass::analyzeScalarVToCDependencies( bool hasVectorDep = hasVectorOpInDefChain(bound, visited, &handledScalarValues); if (!hasVectorDep) { - // The trace may have been stopped by an already-handled scalar - // (e.g. a tensor.extract result transferred by the extract pass). - // In that case the bound itself needs no new transfer, but the - // enclosing loop still becomes "handled" so that ifOp conditions - // inside it are not redundantly detected. + // Chain is vector-only but was stopped by an already-handled scalar: + // no new dep, but mark the loop handled to suppress inner ifOp deps. llvm::DenseSet visitedNoStop; if (!hasVectorOpInDefChain(bound, visitedNoStop)) { continue; @@ -1130,9 +1108,8 @@ void DataDependencyAnalysisPass::analyzeScalarVToCDependencies( LOG_DEBUG("Found scalar V->C dependency from forOp bounds: " << bound << "\n"); - // Mark this for-loop as handled so that ifOp conditions nested inside - // it are skipped — they will be computable on CUBE via the induction - // variable once the bounds are transferred. + // Mark the loop handled: inner ifOp conditions are then computable on + // CUBE via the induction variable. handledForOps.insert(forOp); // Record a single V->C dependency using the for-loop op's own block_id @@ -1161,9 +1138,8 @@ void DataDependencyAnalysisPass::analyzeScalarVToCDependencies( return; } - // Skip if this ifOp is nested inside a for-loop whose bounds have already - // been transferred. The induction variable carries the transferred values, - // so any scalar condition computed from it is already available on CUBE. + // Skip if nested in a handled loop: its condition is then computable on + // CUBE via the induction variable. mlir::Operation *parent = ifOp->getParentOp(); while (parent) { if (auto parentForOp = dyn_cast(parent)) { diff --git a/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/InterCoreTransferAndSync.cpp b/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/InterCoreTransferAndSync.cpp index db4f9d2266..c3571b000c 100644 --- a/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/InterCoreTransferAndSync.cpp +++ b/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/InterCoreTransferAndSync.cpp @@ -547,9 +547,7 @@ Operation *InterCoreTransferAndSyncPass::insertVectorToCubeTransfer( int cubeBlockId = CVPipeline::getOpBlockId(cubeStartOp).value_or(-1); if (isScalarDependency(dep.value)) { - // Insert the store right after srcValue's defining op (when it has one) - // so that the store dominates the load and the scalar consumers. Falling - // back to vectorEndOp keeps behavior for block-argument / external values. + // Place the store right after the extract so it dominates the load. mlir::Operation *srcDefOp = srcValue.getDefiningOp(); if (srcDefOp) { builder.setInsertionPointAfter(srcDefOp); @@ -576,9 +574,8 @@ Operation *InterCoreTransferAndSyncPass::insertVectorToCubeTransfer( attachCrossCoreDeps(sendOp, transferIndex, CVPipeline::crossCoreProducerId, builder); LOG_DEBUG("before readFromSSBuffer\n"); - // Place the load after the store when both are in the same MLIR block, - // otherwise the load would read an uninitialized SSBuffer slot (store and - // load can share a block when producer block == consumer block). + // When store and load share a block (producer == consumer block), load + // after the store to avoid reading an uninitialized slot. if (sendOp && cubeStartOp && sendOp->getBlock() == cubeStartOp->getBlock() && !sendOp->isBeforeInBlock(cubeStartOp)) { @@ -672,9 +669,8 @@ Operation *InterCoreTransferAndSyncPass::insertVectorToCubeTransfer( } for (Operation *user : users) { LOG_DEBUG("[v->c user]" << *user << "\n"); - // Do not rewrite the store/send op itself: it must keep referencing the - // original srcValue, otherwise the value stored into SSBuffer would be the - // just-loaded receiveValue (a store→load self loop). + // Keep the store referencing srcValue, else it stores the loaded value + // back (store→load self loop). if (user == sendOp) { continue; } @@ -812,9 +808,7 @@ InterCoreTransferAndSyncPass::getTransferPipeConfig(Operation *transferOp, config.srcCoreType = "VECTOR"; config.dstCoreType = "CUBE"; } else if (isa(transferOp)) { - // Scalar transfers use PIPE_S (scalar pipe). Using PIPE_V/PIPE_FIX for - // these syncs shares flag space with vector/fix tensor transfers and can - // deadlock the cores; PIPE_S keeps the scalar sync isolated. + // Scalar sync uses PIPE_S to stay isolated from tensor flag space. config.forReadTPipe = pipeSAttr; config.forReadPipe = pipeSAttr; config.forWriteTPipe = pipeSAttr; diff --git a/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/SeparateCVScope.cpp b/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/SeparateCVScope.cpp index d5a366160a..a36369d7c8 100644 --- a/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/SeparateCVScope.cpp +++ b/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/SeparateCVScope.cpp @@ -1014,12 +1014,8 @@ void mlir::triton::SeparateCVScopePass::runOnOperation() { UnitAttr::get(scopeOp->getContext())); }); - // Eliminate redundant store→load pairs in VECTOR scopes. - // InterCoreTransferAndSync inserts llvm.store (send to SSBuffer) followed by - // llvm.load (receive on the consumer side). After scope separation the - // VECTOR scope contains both: the store sends the value, and the load re-reads - // it for use as a for-loop bound. The load is redundant — the stored value - // is already live — so replace loaded values with the stored value. + // In VECTOR scopes the SSBuffer store is followed by a redundant load (re-read + // for a for-loop bound): replace the load with the stored value and erase it. module.walk([](scope::ScopeOp scopeOp) { auto coreTypeAttr = scopeOp->getAttrOfType(hivm::TCoreTypeAttr::name); From 9a7e1048acccf8c18ba6965070dd9c0ae2954def Mon Sep 17 00:00:00 2001 From: 1801ZDL <1241272204@qq.com> Date: Thu, 13 Aug 2026 09:54:18 +0800 Subject: [PATCH 7/9] fix: remove VECTOR redundant-load cleanup (mechanism A) for E2E test Remove the store->load cleanup in VECTOR scopes to test whether downstream dead-code elimination handles the redundant VECTOR-side load. Re-add if the E2E test shows it is still required. Co-Authored-By: DeLong code --- .../SplitDataflow/SeparateCVScope.cpp | 50 +------------------ 1 file changed, 2 insertions(+), 48 deletions(-) diff --git a/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/SeparateCVScope.cpp b/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/SeparateCVScope.cpp index a36369d7c8..0184035b74 100644 --- a/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/SeparateCVScope.cpp +++ b/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/SeparateCVScope.cpp @@ -1014,54 +1014,8 @@ void mlir::triton::SeparateCVScopePass::runOnOperation() { UnitAttr::get(scopeOp->getContext())); }); - // In VECTOR scopes the SSBuffer store is followed by a redundant load (re-read - // for a for-loop bound): replace the load with the stored value and erase it. - module.walk([](scope::ScopeOp scopeOp) { - auto coreTypeAttr = - scopeOp->getAttrOfType(hivm::TCoreTypeAttr::name); - if (!coreTypeAttr || coreTypeAttr.getTcoretype() != hivm::TCoreType::VECTOR) { - return; - } - - llvm::DenseMap storedValues; - scopeOp.walk([&](memref::StoreOp storeOp) { - auto transferIdAttr = - storeOp->getAttrOfType(CVPipeline::kTransferId); - if (!transferIdAttr) { - return; - } - int64_t tid = transferIdAttr.getInt(); - storedValues[tid] = storeOp.getValue(); - }); - - if (storedValues.empty()) { - return; - } - - llvm::SmallVector deadLoads; - scopeOp.walk([&](memref::LoadOp loadOp) { - auto transferIdAttr = - loadOp->getAttrOfType(CVPipeline::kTransferId); - if (!transferIdAttr) { - return; - } - int64_t tid = transferIdAttr.getInt(); - auto it = storedValues.find(tid); - if (it == storedValues.end()) { - return; - } - mlir::Value storeVal = it->second; - if (storeVal == loadOp.getResult()) { - return; - } - loadOp.replaceAllUsesWith(storeVal); - deadLoads.push_back(loadOp); - }); - for (memref::LoadOp loadOp : deadLoads) { - loadOp->erase(); - } - }); - + // TEMP: mechanism A (VECTOR redundant-load cleanup) removed to test whether + // downstream dead-code elimination handles the redundant VECTOR-side load. debugDumpOperation("after SeparateCVScopePass", module.getOperation()); } From ccc26a3acecdafb6ba997086f5eb374a584de2c3 Mon Sep 17 00:00:00 2001 From: 1801ZDL <1241272204@qq.com> Date: Thu, 13 Aug 2026 10:14:18 +0800 Subject: [PATCH 8/9] fix: restore VECTOR redundant-load cleanup (mechanism A) E2E confirmed the VECTOR-side store->load is a redundant operation (re-reads the value the VECTOR core just stored for its own loop bound), not dead code: downstream DCE does not remove it because it feeds the loop bound. Restore the separate-cv-scope cleanup that replaces the load with the stored value. Co-Authored-By: DeLong code --- .../SplitDataflow/SeparateCVScope.cpp | 50 ++++++++++++++++++- 1 file changed, 48 insertions(+), 2 deletions(-) diff --git a/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/SeparateCVScope.cpp b/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/SeparateCVScope.cpp index 0184035b74..a36369d7c8 100644 --- a/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/SeparateCVScope.cpp +++ b/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/SeparateCVScope.cpp @@ -1014,8 +1014,54 @@ void mlir::triton::SeparateCVScopePass::runOnOperation() { UnitAttr::get(scopeOp->getContext())); }); - // TEMP: mechanism A (VECTOR redundant-load cleanup) removed to test whether - // downstream dead-code elimination handles the redundant VECTOR-side load. + // In VECTOR scopes the SSBuffer store is followed by a redundant load (re-read + // for a for-loop bound): replace the load with the stored value and erase it. + module.walk([](scope::ScopeOp scopeOp) { + auto coreTypeAttr = + scopeOp->getAttrOfType(hivm::TCoreTypeAttr::name); + if (!coreTypeAttr || coreTypeAttr.getTcoretype() != hivm::TCoreType::VECTOR) { + return; + } + + llvm::DenseMap storedValues; + scopeOp.walk([&](memref::StoreOp storeOp) { + auto transferIdAttr = + storeOp->getAttrOfType(CVPipeline::kTransferId); + if (!transferIdAttr) { + return; + } + int64_t tid = transferIdAttr.getInt(); + storedValues[tid] = storeOp.getValue(); + }); + + if (storedValues.empty()) { + return; + } + + llvm::SmallVector deadLoads; + scopeOp.walk([&](memref::LoadOp loadOp) { + auto transferIdAttr = + loadOp->getAttrOfType(CVPipeline::kTransferId); + if (!transferIdAttr) { + return; + } + int64_t tid = transferIdAttr.getInt(); + auto it = storedValues.find(tid); + if (it == storedValues.end()) { + return; + } + mlir::Value storeVal = it->second; + if (storeVal == loadOp.getResult()) { + return; + } + loadOp.replaceAllUsesWith(storeVal); + deadLoads.push_back(loadOp); + }); + for (memref::LoadOp loadOp : deadLoads) { + loadOp->erase(); + } + }); + debugDumpOperation("after SeparateCVScopePass", module.getOperation()); } From ba58e929794ff0ebb77694da919b6150661462e2 Mon Sep 17 00:00:00 2001 From: 1801ZDL <1241272204@qq.com> Date: Thu, 13 Aug 2026 12:32:30 +0800 Subject: [PATCH 9/9] refactor: merge for-bound and if-condition scalar detection Remove the handledForOps cross-suppression between for-loop bounds and nested if conditions: whether a loop bound has a scalar dep does not determine whether an if condition inside it does (and vice versa). The outer-loop skip was an over-approximation that could miss a real if condition derived from an independent vector-only chain. Merge both into analyzeScalarControlFlowDependencies: each bound/condition is checked independently against its own defining chain, suppressed only by the shared handledScalarValues stop-set (already-transferred scalars) or by not reaching a vector-only op. handledForOps is removed. Co-Authored-By: DeLong code --- .../SplitDataflow/DataDependencyAnalysis.h | 6 +- .../SplitDataflow/DataDependencyAnalysis.cpp | 72 +++++++------------ 2 files changed, 30 insertions(+), 48 deletions(-) diff --git a/third_party/ascend/include/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.h b/third_party/ascend/include/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.h index f6c21d8dbe..e3108d5a84 100644 --- a/third_party/ascend/include/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.h +++ b/third_party/ascend/include/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.h @@ -147,8 +147,10 @@ class DataDependencyAnalysisPass void analyzeScalarVToCDependencies(DataDependencyInfo &info); void analyzeScalarExtractDependencies( DataDependencyInfo &info, - llvm::DenseSet &handledScalarValues, - llvm::DenseSet &handledForOps); + llvm::DenseSet &handledScalarValues); + void analyzeScalarControlFlowDependencies( + DataDependencyInfo &info, + llvm::DenseSet &handledScalarValues); void analyzeMemoryEffect(DataDependencyInfo &info); std::pair findCommonLevelBlockIds(DataDependencyInfo &info, diff --git a/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.cpp b/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.cpp index 6d2ac3a825..e9dd552aba 100644 --- a/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.cpp +++ b/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.cpp @@ -921,8 +921,7 @@ static bool ifOpHasCubeOps( // Detect scalars extracted from VECTOR-only tensors and consumed by CUBE. void DataDependencyAnalysisPass::analyzeScalarExtractDependencies( DataDependencyInfo &info, - llvm::DenseSet &handledScalarValues, - llvm::DenseSet &handledForOps) { + llvm::DenseSet &handledScalarValues) { auto &blockInfoMap = info.getBlockInfoMap(); auto &v2cDependencies = info.getV2CDependencies(); @@ -1023,16 +1022,6 @@ void DataDependencyAnalysisPass::analyzeScalarExtractDependencies( LOG_DEBUG("Found scalar V->C dependency from tensor.extract: " << scalarResult << "\n"); handledScalarValues.insert(scalarResult); - // Enclosing loop's bounds/conditions derive from this scalar: no further - // transfers needed for them. - mlir::Operation *parent = extractOp->getParentOp(); - while (parent) { - if (auto parentForOp = dyn_cast(parent)) { - handledForOps.insert(parentForOp); - break; - } - parent = parent->getParentOp(); - } } }); } @@ -1048,17 +1037,31 @@ void DataDependencyAnalysisPass::analyzeScalarVToCDependencies( // Scalars already transferred; downstream derivatives need no new dep. llvm::DenseSet handledScalarValues; - // For-loops whose bounds were transferred: inner ifOp conditions are then - // computable via the induction variable, so no separate deps are needed. - llvm::DenseSet handledForOps; - LOG_DEBUG("Analyzing scalar V->C dependencies from control flow ops...\n"); - // ---- tensor.extract ops ---- // Detect scalars extracted from VECTOR-produced tensors and consumed by CUBE // blocks. The extract result is the natural scalar dependency boundary: // transferring it via SSBuffer avoids the need for 1-D tensor CopyOps. - analyzeScalarExtractDependencies(info, handledScalarValues, handledForOps); + analyzeScalarExtractDependencies(info, handledScalarValues); + + // For-loop bounds and if conditions: each scalar is checked independently + // against its own defining chain; there is no cross-suppression between a + // loop and an if nested inside it. + analyzeScalarControlFlowDependencies(info, handledScalarValues); + + LOG_DEBUG("Scalar V->C dependency analysis complete.\n"); +} + +// Detect scalar V->C deps from scf.for loop bounds and scf.if conditions. +// For-loops must contain CUBE ops; bounds/conditions must be scalars whose +// defining chain reaches a vector-only op. For/if are checked independently: +// a handled loop bound does not suppress an if condition inside it, and vice +// versa. +void DataDependencyAnalysisPass::analyzeScalarControlFlowDependencies( + DataDependencyInfo &info, + llvm::DenseSet &handledScalarValues) { + auto &blockInfoMap = info.getBlockInfoMap(); + auto &v2cDependencies = info.getV2CDependencies(); // ---- scf.for loop bounds ---- module.walk([&](scf::ForOp forOp) { @@ -1086,16 +1089,9 @@ void DataDependencyAnalysisPass::analyzeScalarVToCDependencies( } llvm::DenseSet visited; - bool hasVectorDep = - hasVectorOpInDefChain(bound, visited, &handledScalarValues); - if (!hasVectorDep) { - // Chain is vector-only but was stopped by an already-handled scalar: - // no new dep, but mark the loop handled to suppress inner ifOp deps. - llvm::DenseSet visitedNoStop; - if (!hasVectorOpInDefChain(bound, visitedNoStop)) { - continue; - } - handledForOps.insert(forOp); + if (!hasVectorOpInDefChain(bound, visited, &handledScalarValues)) { + // Bound is not from a vector-only chain, or is already covered by a + // transferred scalar (stop-set): no new dep needed. continue; } @@ -1108,10 +1104,6 @@ void DataDependencyAnalysisPass::analyzeScalarVToCDependencies( LOG_DEBUG("Found scalar V->C dependency from forOp bounds: " << bound << "\n"); - // Mark the loop handled: inner ifOp conditions are then computable on - // CUBE via the induction variable. - handledForOps.insert(forOp); - // Record a single V->C dependency using the for-loop op's own block_id // as the consumer. The bound is also used by ops inside the loop body, // but those are in the same CUBE scope after CV separation and will use @@ -1138,18 +1130,6 @@ void DataDependencyAnalysisPass::analyzeScalarVToCDependencies( return; } - // Skip if nested in a handled loop: its condition is then computable on - // CUBE via the induction variable. - mlir::Operation *parent = ifOp->getParentOp(); - while (parent) { - if (auto parentForOp = dyn_cast(parent)) { - if (handledForOps.contains(parentForOp)) { - return; - } - } - parent = parent->getParentOp(); - } - mlir::Value condition = ifOp.getCondition(); if (!isa(condition.getType())) { return; @@ -1162,6 +1142,8 @@ void DataDependencyAnalysisPass::analyzeScalarVToCDependencies( llvm::DenseSet visited; if (!hasVectorOpInDefChain(condition, visited, &handledScalarValues)) { + // Condition is not from a vector-only chain, or is already covered by a + // transferred scalar (stop-set): no new dep needed. return; } @@ -1188,8 +1170,6 @@ void DataDependencyAnalysisPass::analyzeScalarVToCDependencies( producerId, consumerId, info); handledScalarValues.insert(condition); }); - - LOG_DEBUG("Scalar V->C dependency analysis complete.\n"); } void DataDependencyAnalysisPass::analyzeMemoryEffect(DataDependencyInfo &info) {