diff --git a/mlir/include/mlir/Transforms/DialectConversion.h b/mlir/include/mlir/Transforms/DialectConversion.h index 5f680e8eca755..60113bdef16a2 100644 --- a/mlir/include/mlir/Transforms/DialectConversion.h +++ b/mlir/include/mlir/Transforms/DialectConversion.h @@ -1124,17 +1124,6 @@ struct ConversionConfig { // already been modified) and iterators into past IR state cannot be // represented at the moment. RewriterBase::Listener *listener = nullptr; - - /// If set to "true", the dialect conversion attempts to build source/target/ - /// argument materializations through the type converter API in lieu of - /// builtin.unrealized_conversion_cast ops. The conversion process fails if - /// at least one materialization could not be built. - /// - /// If set to "false", the dialect conversion does not does not build any - /// custom materializations and instead inserts - /// builtin.unrealized_conversion_cast ops to ensure that the resulting IR - /// is valid. - bool buildMaterializations = true; }; //===----------------------------------------------------------------------===// diff --git a/mlir/lib/Transforms/Utils/DialectConversion.cpp b/mlir/lib/Transforms/Utils/DialectConversion.cpp index cc9c9495e5155..b23fb97959ed6 100644 --- a/mlir/lib/Transforms/Utils/DialectConversion.cpp +++ b/mlir/lib/Transforms/Utils/DialectConversion.cpp @@ -702,12 +702,14 @@ class UnresolvedMaterializationRewrite : public OperationRewrite { return rewrite->getKind() == Kind::UnresolvedMaterialization; } - void rollback() override; - UnrealizedConversionCastOp getOperation() const { return cast(op); } + void rollback() override; + + void cleanup(RewriterBase &rewriter) override; + /// Return the type converter of this materialization (which may be null). const TypeConverter *getConverter() const { return converterAndKind.getPointer(); @@ -764,7 +766,7 @@ namespace detail { struct ConversionPatternRewriterImpl : public RewriterBase::Listener { explicit ConversionPatternRewriterImpl(MLIRContext *ctx, const ConversionConfig &config) - : context(ctx), eraseRewriter(ctx), config(config) {} + : context(ctx), config(config) {} //===--------------------------------------------------------------------===// // State Management @@ -832,7 +834,6 @@ struct ConversionPatternRewriterImpl : public RewriterBase::Listener { //===--------------------------------------------------------------------===// // Materializations //===--------------------------------------------------------------------===// - /// Build an unresolved materialization operation given an output type and set /// of input operands. Value buildUnresolvedMaterialization(MaterializationKind kind, @@ -881,7 +882,7 @@ struct ConversionPatternRewriterImpl : public RewriterBase::Listener { /// Erase the given op (unless it was already erased). void eraseOp(Operation *op) override { - if (wasErased(op)) + if (erased.contains(op)) return; op->dropAllUses(); RewriterBase::eraseOp(op); @@ -889,24 +890,17 @@ struct ConversionPatternRewriterImpl : public RewriterBase::Listener { /// Erase the given block (unless it was already erased). void eraseBlock(Block *block) override { - if (wasErased(block)) + if (erased.contains(block)) return; assert(block->empty() && "expected empty block"); block->dropAllDefinedValueUses(); RewriterBase::eraseBlock(block); } - bool wasErased(void *ptr) const { return erased.contains(ptr); } - - bool wasErased(OperationRewrite *rewrite) const { - return wasErased(rewrite->getOperation()); - } - void notifyOperationErased(Operation *op) override { erased.insert(op); } void notifyBlockErased(Block *block) override { erased.insert(block); } - private: /// Pointers to all erased operations and blocks. DenseSet erased; }; @@ -918,11 +912,6 @@ struct ConversionPatternRewriterImpl : public RewriterBase::Listener { /// MLIR context. MLIRContext *context; - /// A rewriter that keeps track of ops/block that were already erased and - /// skips duplicate op/block erasures. This rewriter is used during the - /// "cleanup" phase. - SingleEraseRewriter eraseRewriter; - // Mapping between replaced values that differ in type. This happens when // replacing a value with one of a different type. ConversionValueMapping mapping; @@ -1069,6 +1058,10 @@ void UnresolvedMaterializationRewrite::rollback() { op->erase(); } +void UnresolvedMaterializationRewrite::cleanup(RewriterBase &rewriter) { + rewriter.eraseOp(op); +} + void ConversionPatternRewriterImpl::applyRewrites() { // Commit all rewrites. IRRewriter rewriter(context, config.listener); @@ -1076,6 +1069,7 @@ void ConversionPatternRewriterImpl::applyRewrites() { rewrite->commit(rewriter); // Clean up all rewrites. + SingleEraseRewriter eraseRewriter(context); for (auto &rewrite : rewrites) rewrite->cleanup(eraseRewriter); } @@ -2359,6 +2353,12 @@ struct OperationConverter { ConversionPatternRewriterImpl &rewriterImpl, DenseMap> &inverseMapping); + /// Legalize any unresolved type materializations. + LogicalResult legalizeUnresolvedMaterializations( + ConversionPatternRewriter &rewriter, + ConversionPatternRewriterImpl &rewriterImpl, + DenseMap> &inverseMapping); + /// Legalize an operation result that was marked as "erased". LogicalResult legalizeErasedResult(Operation *op, OpResult result, @@ -2405,56 +2405,6 @@ LogicalResult OperationConverter::convert(ConversionPatternRewriter &rewriter, return success(); } -static LogicalResult -legalizeUnresolvedMaterialization(RewriterBase &rewriter, - UnresolvedMaterializationRewrite *rewrite) { - UnrealizedConversionCastOp op = rewrite->getOperation(); - assert(!op.use_empty() && - "expected that dead materializations have already been DCE'd"); - Operation::operand_range inputOperands = op.getOperands(); - Type outputType = op.getResultTypes()[0]; - - // Try to materialize the conversion. - if (const TypeConverter *converter = rewrite->getConverter()) { - rewriter.setInsertionPoint(op); - Value newMaterialization; - switch (rewrite->getMaterializationKind()) { - case MaterializationKind::Argument: - // Try to materialize an argument conversion. - newMaterialization = converter->materializeArgumentConversion( - rewriter, op->getLoc(), outputType, inputOperands); - if (newMaterialization) - break; - // If an argument materialization failed, fallback to trying a target - // materialization. - [[fallthrough]]; - case MaterializationKind::Target: - newMaterialization = converter->materializeTargetConversion( - rewriter, op->getLoc(), outputType, inputOperands); - break; - case MaterializationKind::Source: - newMaterialization = converter->materializeSourceConversion( - rewriter, op->getLoc(), outputType, inputOperands); - break; - } - if (newMaterialization) { - assert(newMaterialization.getType() == outputType && - "materialization callback produced value of incorrect type"); - rewriter.replaceOp(op, newMaterialization); - return success(); - } - } - - InFlightDiagnostic diag = op->emitError() - << "failed to legalize unresolved materialization " - "from (" - << inputOperands.getTypes() << ") to " << outputType - << " that remained live after conversion"; - diag.attachNote(op->getUsers().begin()->getLoc()) - << "see existing live user here: " << *op->getUsers().begin(); - return failure(); -} - LogicalResult OperationConverter::convertOperations(ArrayRef ops) { if (ops.empty()) return success(); @@ -2496,37 +2446,6 @@ LogicalResult OperationConverter::convertOperations(ArrayRef ops) { } else { rewriterImpl.applyRewrites(); } - - // Gather all unresolved materializations. - SmallVector allCastOps; - DenseMap rewriteMap; - for (std::unique_ptr &rewrite : rewriterImpl.rewrites) { - auto *mat = dyn_cast(rewrite.get()); - if (!mat) - continue; - if (rewriterImpl.eraseRewriter.wasErased(mat)) - continue; - allCastOps.push_back(mat->getOperation()); - rewriteMap[mat->getOperation()] = mat; - } - - // Reconcile all UnrealizedConversionCastOps that were inserted by the - // dialect conversion frameworks. (Not the one that were inserted by - // patterns.) - SmallVector remainingCastOps; - reconcileUnrealizedCasts(allCastOps, &remainingCastOps); - - // Try to legalize all unresolved materializations. - if (config.buildMaterializations) { - IRRewriter rewriter(rewriterImpl.context, config.listener); - for (UnrealizedConversionCastOp castOp : remainingCastOps) { - auto it = rewriteMap.find(castOp.getOperation()); - assert(it != rewriteMap.end() && "inconsistent state"); - if (failed(legalizeUnresolvedMaterialization(rewriter, it->second))) - return failure(); - } - } - return success(); } @@ -2540,6 +2459,9 @@ OperationConverter::finalize(ConversionPatternRewriter &rewriter) { if (failed(legalizeConvertedOpResultTypes(rewriter, rewriterImpl, inverseMapping))) return failure(); + if (failed(legalizeUnresolvedMaterializations(rewriter, rewriterImpl, + inverseMapping))) + return failure(); return success(); } @@ -2655,6 +2577,279 @@ LogicalResult OperationConverter::legalizeConvertedArgumentTypes( return success(); } +/// Replace the results of a materialization operation with the given values. +static void +replaceMaterialization(ConversionPatternRewriterImpl &rewriterImpl, + ResultRange matResults, ValueRange values, + DenseMap> &inverseMapping) { + matResults.replaceAllUsesWith(values); + + // For each of the materialization results, update the inverse mappings to + // point to the replacement values. + for (auto [matResult, newValue] : llvm::zip(matResults, values)) { + auto inverseMapIt = inverseMapping.find(matResult); + if (inverseMapIt == inverseMapping.end()) + continue; + + // Update the reverse mapping, or remove the mapping if we couldn't update + // it. Not being able to update signals that the mapping would have become + // circular (i.e. %foo -> newValue -> %foo), which may occur as values are + // propagated through temporary materializations. We simply drop the + // mapping, and let the post-conversion replacement logic handle updating + // uses. + for (Value inverseMapVal : inverseMapIt->second) + if (!rewriterImpl.mapping.tryMap(inverseMapVal, newValue)) + rewriterImpl.mapping.erase(inverseMapVal); + } +} + +/// Compute all of the unresolved materializations that will persist beyond the +/// conversion process, and require inserting a proper user materialization for. +static void computeNecessaryMaterializations( + DenseMap + &materializationOps, + ConversionPatternRewriter &rewriter, + ConversionPatternRewriterImpl &rewriterImpl, + DenseMap> &inverseMapping, + SetVector &necessaryMaterializations) { + // Helper function to check if the given value or a not yet materialized + // replacement of the given value is live. + // Note: `inverseMapping` maps from replaced values to original values. + auto isLive = [&](Value value) { + auto findFn = [&](Operation *user) { + auto matIt = materializationOps.find(user); + if (matIt != materializationOps.end()) + return !necessaryMaterializations.count(matIt->second); + return rewriterImpl.isOpIgnored(user); + }; + // A worklist is needed because a value may have gone through a chain of + // replacements and each of the replaced values may have live users. + SmallVector worklist; + worklist.push_back(value); + while (!worklist.empty()) { + Value next = worklist.pop_back_val(); + if (llvm::find_if_not(next.getUsers(), findFn) != next.user_end()) + return true; + // This value may be replacing another value that has a live user. + llvm::append_range(worklist, inverseMapping.lookup(next)); + } + return false; + }; + + llvm::unique_function lookupRemappedValue = + [&](Value invalidRoot, Value value, Type type) { + // Check to see if the input operation was remapped to a variant of the + // output. + Value remappedValue = rewriterImpl.mapping.lookupOrDefault(value, type); + if (remappedValue.getType() == type && remappedValue != invalidRoot) + return remappedValue; + + // Check to see if the input is a materialization operation that + // provides an inverse conversion. We just check blindly for + // UnrealizedConversionCastOp here, but it has no effect on correctness. + auto inputCastOp = value.getDefiningOp(); + if (inputCastOp && inputCastOp->getNumOperands() == 1) + return lookupRemappedValue(invalidRoot, inputCastOp->getOperand(0), + type); + + return Value(); + }; + + SetVector worklist; + for (auto &rewrite : rewriterImpl.rewrites) { + auto *mat = dyn_cast(rewrite.get()); + if (!mat) + continue; + materializationOps.try_emplace(mat->getOperation(), mat); + worklist.insert(mat); + } + while (!worklist.empty()) { + UnresolvedMaterializationRewrite *mat = worklist.pop_back_val(); + UnrealizedConversionCastOp op = mat->getOperation(); + + // We currently only handle target materializations here. + assert(op->getNumResults() == 1 && "unexpected materialization type"); + OpResult opResult = op->getOpResult(0); + Type outputType = opResult.getType(); + Operation::operand_range inputOperands = op.getOperands(); + + // Try to forward propagate operands for user conversion casts that result + // in the input types of the current cast. + for (Operation *user : llvm::make_early_inc_range(opResult.getUsers())) { + auto castOp = dyn_cast(user); + if (!castOp) + continue; + if (castOp->getResultTypes() == inputOperands.getTypes()) { + replaceMaterialization(rewriterImpl, user->getResults(), inputOperands, + inverseMapping); + necessaryMaterializations.remove(materializationOps.lookup(user)); + } + } + + // Try to avoid materializing a resolved materialization if possible. + // Handle the case of a 1-1 materialization. + if (inputOperands.size() == 1) { + // Check to see if the input operation was remapped to a variant of the + // output. + Value remappedValue = + lookupRemappedValue(opResult, inputOperands[0], outputType); + if (remappedValue && remappedValue != opResult) { + replaceMaterialization(rewriterImpl, opResult, remappedValue, + inverseMapping); + necessaryMaterializations.remove(mat); + continue; + } + } else { + // TODO: Avoid materializing other types of conversions here. + } + + // If the materialization does not have any live users, we don't need to + // generate a user materialization for it. + bool isMaterializationLive = isLive(opResult); + if (!isMaterializationLive) + continue; + if (!necessaryMaterializations.insert(mat)) + continue; + + // Reprocess input materializations to see if they have an updated status. + for (Value input : inputOperands) { + if (auto parentOp = input.getDefiningOp()) { + if (auto *mat = materializationOps.lookup(parentOp)) + worklist.insert(mat); + } + } + } +} + +/// Legalize the given unresolved materialization. Returns success if the +/// materialization was legalized, failure otherise. +static LogicalResult legalizeUnresolvedMaterialization( + UnresolvedMaterializationRewrite &mat, + DenseMap + &materializationOps, + ConversionPatternRewriter &rewriter, + ConversionPatternRewriterImpl &rewriterImpl, + DenseMap> &inverseMapping) { + auto findLiveUser = [&](auto &&users) { + auto liveUserIt = llvm::find_if_not( + users, [&](Operation *user) { return rewriterImpl.isOpIgnored(user); }); + return liveUserIt == users.end() ? nullptr : *liveUserIt; + }; + + llvm::unique_function lookupRemappedValue = + [&](Value value, Type type) { + // Check to see if the input operation was remapped to a variant of the + // output. + Value remappedValue = rewriterImpl.mapping.lookupOrDefault(value, type); + if (remappedValue.getType() == type) + return remappedValue; + return Value(); + }; + + UnrealizedConversionCastOp op = mat.getOperation(); + if (!rewriterImpl.ignoredOps.insert(op)) + return success(); + + // We currently only handle target materializations here. + OpResult opResult = op->getOpResult(0); + Operation::operand_range inputOperands = op.getOperands(); + Type outputType = opResult.getType(); + + // If any input to this materialization is another materialization, resolve + // the input first. + for (Value value : op->getOperands()) { + auto valueCast = value.getDefiningOp(); + if (!valueCast) + continue; + + auto matIt = materializationOps.find(valueCast); + if (matIt != materializationOps.end()) + if (failed(legalizeUnresolvedMaterialization( + *matIt->second, materializationOps, rewriter, rewriterImpl, + inverseMapping))) + return failure(); + } + + // Perform a last ditch attempt to avoid materializing a resolved + // materialization if possible. + // Handle the case of a 1-1 materialization. + if (inputOperands.size() == 1) { + // Check to see if the input operation was remapped to a variant of the + // output. + Value remappedValue = lookupRemappedValue(inputOperands[0], outputType); + if (remappedValue && remappedValue != opResult) { + replaceMaterialization(rewriterImpl, opResult, remappedValue, + inverseMapping); + return success(); + } + } else { + // TODO: Avoid materializing other types of conversions here. + } + + // Try to materialize the conversion. + if (const TypeConverter *converter = mat.getConverter()) { + rewriter.setInsertionPoint(op); + Value newMaterialization; + switch (mat.getMaterializationKind()) { + case MaterializationKind::Argument: + // Try to materialize an argument conversion. + newMaterialization = converter->materializeArgumentConversion( + rewriter, op->getLoc(), outputType, inputOperands); + if (newMaterialization) + break; + // If an argument materialization failed, fallback to trying a target + // materialization. + [[fallthrough]]; + case MaterializationKind::Target: + newMaterialization = converter->materializeTargetConversion( + rewriter, op->getLoc(), outputType, inputOperands); + break; + case MaterializationKind::Source: + newMaterialization = converter->materializeSourceConversion( + rewriter, op->getLoc(), outputType, inputOperands); + break; + } + if (newMaterialization) { + assert(newMaterialization.getType() == outputType && + "materialization callback produced value of incorrect type"); + replaceMaterialization(rewriterImpl, opResult, newMaterialization, + inverseMapping); + return success(); + } + } + + InFlightDiagnostic diag = op->emitError() + << "failed to legalize unresolved materialization " + "from (" + << inputOperands.getTypes() << ") to " << outputType + << " that remained live after conversion"; + if (Operation *liveUser = findLiveUser(op->getUsers())) { + diag.attachNote(liveUser->getLoc()) + << "see existing live user here: " << *liveUser; + } + return failure(); +} + +LogicalResult OperationConverter::legalizeUnresolvedMaterializations( + ConversionPatternRewriter &rewriter, + ConversionPatternRewriterImpl &rewriterImpl, + DenseMap> &inverseMapping) { + // As an initial step, compute all of the inserted materializations that we + // expect to persist beyond the conversion process. + DenseMap materializationOps; + SetVector necessaryMaterializations; + computeNecessaryMaterializations(materializationOps, rewriter, rewriterImpl, + inverseMapping, necessaryMaterializations); + + // Once computed, legalize any necessary materializations. + for (auto *mat : necessaryMaterializations) { + if (failed(legalizeUnresolvedMaterialization( + *mat, materializationOps, rewriter, rewriterImpl, inverseMapping))) + return failure(); + } + return success(); +} + LogicalResult OperationConverter::legalizeErasedResult( Operation *op, OpResult result, ConversionPatternRewriterImpl &rewriterImpl) { diff --git a/mlir/test/Conversion/NVGPUToNVVM/nvgpu-to-nvvm.mlir b/mlir/test/Conversion/NVGPUToNVVM/nvgpu-to-nvvm.mlir index 75362378daaaa..156a8a468d5b4 100644 --- a/mlir/test/Conversion/NVGPUToNVVM/nvgpu-to-nvvm.mlir +++ b/mlir/test/Conversion/NVGPUToNVVM/nvgpu-to-nvvm.mlir @@ -1286,6 +1286,7 @@ func.func @warpgroup_matrix_multiply_m128n128k64( // CHECK-DAG: %[[S0:.+]] = builtin.unrealized_conversion_cast %[[arg0]] : !nvgpu.warpgroup.descriptor> to i64 // CHECK-DAG: %[[S1:.+]] = builtin.unrealized_conversion_cast %[[arg1]] : !nvgpu.warpgroup.descriptor> to i64 +// CHECK-DAG: %[[S2:.+]] = builtin.unrealized_conversion_cast %[[arg2]] : memref<128x128xf32, 3> to !llvm.struct<(ptr<3>, ptr<3>, i64, array<2 x i64>, array<2 x i64>)> // CHECK: %[[S3:.+]] = llvm.mlir.constant(0.000000e+00 : f32) : f32 // CHECK: %[[S4:.+]] = llvm.mlir.undef : !llvm.struct<(struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)>, struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)>)> // CHECK: %[[S5:.+]] = llvm.extractvalue %[[S4]][0] : !llvm.struct<(struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)>, struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)>)> @@ -1298,8 +1299,8 @@ func.func @warpgroup_matrix_multiply_m128n128k64( // CHECK: %[[S136:.+]] = llvm.insertvalue %[[S134]], %[[S135]][1] : !llvm.struct<(struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)>, struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)>)> // CHECK: nvvm.wgmma.fence.aligned // CHECK: %[[S137:.+]] = llvm.mlir.undef : !llvm.struct<(struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)>, struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)>)> -// CHECK: %[[S138:.+]] = llvm.extractvalue %{{.*}}[0] : !llvm.struct<(struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)>, struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)>)> -// CHECK: %[[S139:.+]] = nvvm.wgmma.mma_async %[[S0]], %[[S1]], %[[S138]], , D[, , ], A[, , ], B[, , ] : !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)> -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)> +// CHECK: %[[S138:.+]] = llvm.extractvalue %136[0] : !llvm.struct<(struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)>, struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)>)> +// CHECK: %[[S139:.+]] = nvvm.wgmma.mma_async %[[S0]], %1, %[[S138]], , D[, , ], A[, , ], B[, , ] : !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)> -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32, f32)> // CHECK: nvvm.wgmma.mma_async // CHECK: nvvm.wgmma.mma_async // CHECK: %[[S154:.+]] = nvvm.wgmma.mma_async diff --git a/mlir/test/Dialect/Bufferization/Transforms/finalizing-bufferize.mlir b/mlir/test/Dialect/Bufferization/Transforms/finalizing-bufferize.mlir index ab18ce05e355d..a192434c5accf 100644 --- a/mlir/test/Dialect/Bufferization/Transforms/finalizing-bufferize.mlir +++ b/mlir/test/Dialect/Bufferization/Transforms/finalizing-bufferize.mlir @@ -80,7 +80,6 @@ func.func @no_layout_to_dyn_layout_cast(%m: memref) -> memref // expected-error @+1 {{failed to legalize unresolved materialization from ('memref') to 'memref>' that remained live after conversion}} %1 = bufferization.to_memref %0 : memref> - // expected-note @below{{see existing live user here}} return %1 : memref> } diff --git a/mlir/test/Transforms/test-legalize-type-conversion.mlir b/mlir/test/Transforms/test-legalize-type-conversion.mlir index f130adff42f8c..cf2c9f6a8ec44 100644 --- a/mlir/test/Transforms/test-legalize-type-conversion.mlir +++ b/mlir/test/Transforms/test-legalize-type-conversion.mlir @@ -4,7 +4,6 @@ func.func @test_invalid_arg_materialization( // expected-error@below {{failed to legalize unresolved materialization from () to 'i16' that remained live after conversion}} %arg0: i16) { - // expected-note@below{{see existing live user here}} "foo.return"(%arg0) : (i16) -> () } @@ -23,7 +22,6 @@ func.func @test_valid_arg_materialization(%arg0: i64) { func.func @test_invalid_result_materialization() { // expected-error@below {{failed to legalize unresolved materialization from ('f64') to 'f16' that remained live after conversion}} %result = "test.type_producer"() : () -> f16 - // expected-note@below{{see existing live user here}} "foo.return"(%result) : (f16) -> () } @@ -32,7 +30,6 @@ func.func @test_invalid_result_materialization() { func.func @test_invalid_result_materialization() { // expected-error@below {{failed to legalize unresolved materialization from ('f64') to 'f16' that remained live after conversion}} %result = "test.type_producer"() : () -> f16 - // expected-note@below{{see existing live user here}} "foo.return"(%result) : (f16) -> () } @@ -52,7 +49,6 @@ func.func @test_transitive_use_materialization() { func.func @test_transitive_use_invalid_materialization() { // expected-error@below {{failed to legalize unresolved materialization from ('f64') to 'f16' that remained live after conversion}} %result = "test.another_type_producer"() : () -> f16 - // expected-note@below{{see existing live user here}} "foo.return"(%result) : (f16) -> () } @@ -103,9 +99,9 @@ func.func @test_block_argument_not_converted() { func.func @test_signature_conversion_no_converter() { "test.signature_conversion_no_converter"() ({ // expected-error@below {{failed to legalize unresolved materialization from ('f64') to 'f32' that remained live after conversion}} + // expected-note@below {{see existing live user here}} ^bb0(%arg0: f32): "test.type_consumer"(%arg0) : (f32) -> () - // expected-note@below{{see existing live user here}} "test.return"(%arg0) : (f32) -> () }) : () -> () return