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..e3108d5a84 100644 --- a/third_party/ascend/include/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.h +++ b/third_party/ascend/include/DynamicCVPipeline/SplitDataflow/DataDependencyAnalysis.h @@ -144,6 +144,13 @@ class DataDependencyAnalysisPass mlir::Operation *predOp, mlir::Operation *nextOp); void analyzeExternalInputs(DataDependencyInfo &info); void analyzeExternalOutputs(DataDependencyInfo &info); + void analyzeScalarVToCDependencies(DataDependencyInfo &info); + void analyzeScalarExtractDependencies( + DataDependencyInfo &info, + 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/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..a5c383ccb6 100644 --- a/third_party/ascend/lib/DynamicCVPipeline/Common/SSBufferManager.cpp +++ b/third_party/ascend/lib/DynamicCVPipeline/Common/SSBufferManager.cpp @@ -103,8 +103,9 @@ SSBufferManager::writeToSSBuffer(Value value, OpBuilder &builder, int64_t addrValue = addrResult.value(); Location loc = builder.getUnknownLoc(); + // memref.store requires value type == 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 +130,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..dadaa248bc 100644 --- a/third_party/ascend/lib/DynamicCVPipeline/PlanComputeBlock/OpClassifier.cpp +++ b/third_party/ascend/lib/DynamicCVPipeline/PlanComputeBlock/OpClassifier.cpp @@ -703,6 +703,41 @@ void OpClassifierPass::getUpstreamOpsWithMemoryDeps( } } +// arith/math op with a tensor result is VECTOR-only (not 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; +} + +// 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; + 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"); @@ -729,19 +764,20 @@ 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) { + // 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; + } + + // An extract of a VECTOR-only tensor is itself VECTOR. + if (auto extOp = dyn_cast(def)) { + if (hasVectorOnlyProducer(extOp.getTensor())) { + cubeVisited.insert(def); LLVM_DEBUG(DBGS() << "skip " << def->getName().getStringRef() - << ": arith tensor op\n"); + << ": extract of vector-only producer\n"); continue; } } @@ -927,6 +963,16 @@ void OpClassifierPass::propagateCubeUpstreamForOp(Operation *startOp) { continue; if (isa(upstreamOp)) continue; + // 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 tensor is itself VECTOR. + if (auto extOp = dyn_cast(upstreamOp)) { + if (hasVectorOnlyProducer(extOp.getTensor())) + 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..e9dd552aba 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,332 @@ 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; +} + +// Detect scalars extracted from VECTOR-only tensors and consumed by CUBE. +void DataDependencyAnalysisPass::analyzeScalarExtractDependencies( + DataDependencyInfo &info, + llvm::DenseSet &handledScalarValues) { + auto &blockInfoMap = info.getBlockInfoMap(); + auto &v2cDependencies = info.getV2CDependencies(); + + 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; + } + + // 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; + } + + mlir::Value scalarResult = extractOp.getResult(); + if (!isa(scalarResult.getType())) { + return; + } + + // 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; + 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) { + // 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()) { + // One dep per extract; all consumers share a single SSBuffer store. + 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); + } + }); +} + +// 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 transferred; downstream derivatives need no new dep. + llvm::DenseSet handledScalarValues; + + LOG_DEBUG("Analyzing scalar V->C dependencies from control flow ops...\n"); + + // 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); + + // 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) { + 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; + 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; + } + + auto producerIdOpt = CVPipeline::getOpBlockId(defOp); + if (!producerIdOpt) { + continue; + } + int producerId = *producerIdOpt; + + LOG_DEBUG("Found scalar V->C dependency from forOp bounds: " + << bound << "\n"); + + // 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; + } + + 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)) { + // Condition is not from a vector-only chain, or is already covered by a + // transferred scalar (stop-set): no new dep needed. + 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); + }); +} + void DataDependencyAnalysisPass::analyzeMemoryEffect(DataDependencyInfo &info) { auto &memoryDependencies = info.getMemoryDependencies(); LOG_DEBUG("\n=== start mem dep analysis ===\n"); @@ -1055,6 +1386,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..c3571b000c 100644 --- a/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/InterCoreTransferAndSync.cpp +++ b/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/InterCoreTransferAndSync.cpp @@ -547,7 +547,13 @@ Operation *InterCoreTransferAndSyncPass::insertVectorToCubeTransfer( int cubeBlockId = CVPipeline::getOpBlockId(cubeStartOp).value_or(-1); if (isScalarDependency(dep.value)) { - builder.setInsertionPointAfter(vectorEndOp); + // Place the store right after the extract so it dominates the load. + 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 +574,15 @@ Operation *InterCoreTransferAndSyncPass::insertVectorToCubeTransfer( attachCrossCoreDeps(sendOp, transferIndex, CVPipeline::crossCoreProducerId, builder); LOG_DEBUG("before readFromSSBuffer\n"); - builder.setInsertionPoint(cubeStartOp); + // 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)) { + builder.setInsertionPointAfter(sendOp); + } else { + builder.setInsertionPoint(cubeStartOp); + } SmallVector readOps; auto loadedValueOpt = ssbufferManager.readFromSSBuffer(addr, builder, readOps); @@ -655,6 +669,11 @@ Operation *InterCoreTransferAndSyncPass::insertVectorToCubeTransfer( } for (Operation *user : users) { LOG_DEBUG("[v->c user]" << *user << "\n"); + // Keep the store referencing srcValue, else it stores the loaded value + // back (store→load self loop). + if (user == sendOp) { + continue; + } auto userBlockIdOpt = CVPipeline::getOpBlockId(user); if (userBlockIdOpt && *userBlockIdOpt == dep.iniConsumerBlockId) { user->replaceUsesOfWith(srcValue, receiveValue); @@ -789,10 +808,11 @@ 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 sync uses PIPE_S to stay isolated from tensor flag space. + 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..a36369d7c8 100644 --- a/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/SeparateCVScope.cpp +++ b/third_party/ascend/lib/DynamicCVPipeline/SplitDataflow/SeparateCVScope.cpp @@ -22,11 +22,14 @@ #include +#include + #include "llvm/ADT/SmallPtrSet.h" #include "llvm/ADT/SmallVector.h" #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" @@ -38,12 +41,7 @@ #include "mlir/IR/IRMapping.h" #include "mlir/Interfaces/LoopLikeInterface.h" #include "mlir/Pass/Pass.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/Transforms/RegionUtils.h" using namespace mlir; @@ -951,6 +949,7 @@ static void cleanupSsbufferAttrs(Operation *rootOp) { rootOp->walk([](Operation *op) { removeSsbufferAttrs(op); }); } + static LogicalResult separateScopes(func::FuncOp funcOp) { debugDumpOperation("before SeparateCVScope on func", funcOp.getOperation()); @@ -1015,6 +1014,54 @@ 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(); + } + }); + debugDumpOperation("after SeparateCVScopePass", module.getOperation()); }