[Mlir-commits] [mlir] Revert "[mlir][Transforms] Dialect conversion: Make materializations optional" (PR #106778)
Matthias Springer
llvmlistbot at llvm.org
Fri Aug 30 11:57:10 PDT 2024
https://github.com/matthias-springer created https://github.com/llvm/llvm-project/pull/106778
Reverts llvm/llvm-project#104668
This commit triggers an edge case that can cause circular `unrealized_conversion_cast` ops. https://github.com/llvm/llvm-project/pull/106760 may fix it, but it is has other issues. Reverting this PR for now, until I find a solution for that problem.
>From f98550056573efff12b2fb44980f210312b1db8d Mon Sep 17 00:00:00 2001
From: Matthias Springer <me at matthiasspringer.de>
Date: Fri, 30 Aug 2024 11:53:51 -0700
Subject: [PATCH] =?UTF-8?q?Revert=20"[mlir][Transforms]=20Dialect=20conver?=
=?UTF-8?q?sion:=20Make=20materializations=20optional=E2=80=A6"?=
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit
This reverts commit d7073c527457dc0a71126381afb3c6f0efa1821c.
---
.../mlir/Transforms/DialectConversion.h | 11 -
.../Transforms/Utils/DialectConversion.cpp | 393 +++++++++++++-----
.../Conversion/NVGPUToNVVM/nvgpu-to-nvvm.mlir | 5 +-
.../Transforms/finalizing-bufferize.mlir | 1 -
.../test-legalize-type-conversion.mlir | 6 +-
5 files changed, 298 insertions(+), 118 deletions(-)
diff --git a/mlir/include/mlir/Transforms/DialectConversion.h b/mlir/include/mlir/Transforms/DialectConversion.h
index 5f680e8eca7559..60113bdef16a23 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 cc9c9495e5155c..b23fb97959ed67 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<UnrealizedConversionCastOp>(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<void *> 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<Value, SmallVector<Value>> &inverseMapping);
+ /// Legalize any unresolved type materializations.
+ LogicalResult legalizeUnresolvedMaterializations(
+ ConversionPatternRewriter &rewriter,
+ ConversionPatternRewriterImpl &rewriterImpl,
+ DenseMap<Value, SmallVector<Value>> &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<Operation *> ops) {
if (ops.empty())
return success();
@@ -2496,37 +2446,6 @@ LogicalResult OperationConverter::convertOperations(ArrayRef<Operation *> ops) {
} else {
rewriterImpl.applyRewrites();
}
-
- // Gather all unresolved materializations.
- SmallVector<UnrealizedConversionCastOp> allCastOps;
- DenseMap<Operation *, UnresolvedMaterializationRewrite *> rewriteMap;
- for (std::unique_ptr<IRRewrite> &rewrite : rewriterImpl.rewrites) {
- auto *mat = dyn_cast<UnresolvedMaterializationRewrite>(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<UnrealizedConversionCastOp> 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<Value, SmallVector<Value>> &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<Operation *, UnresolvedMaterializationRewrite *>
+ &materializationOps,
+ ConversionPatternRewriter &rewriter,
+ ConversionPatternRewriterImpl &rewriterImpl,
+ DenseMap<Value, SmallVector<Value>> &inverseMapping,
+ SetVector<UnresolvedMaterializationRewrite *> &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<Value> 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<Value(Value, Value, Type)> 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<UnrealizedConversionCastOp>();
+ if (inputCastOp && inputCastOp->getNumOperands() == 1)
+ return lookupRemappedValue(invalidRoot, inputCastOp->getOperand(0),
+ type);
+
+ return Value();
+ };
+
+ SetVector<UnresolvedMaterializationRewrite *> worklist;
+ for (auto &rewrite : rewriterImpl.rewrites) {
+ auto *mat = dyn_cast<UnresolvedMaterializationRewrite>(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<UnrealizedConversionCastOp>(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<UnrealizedConversionCastOp>()) {
+ 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<Operation *, UnresolvedMaterializationRewrite *>
+ &materializationOps,
+ ConversionPatternRewriter &rewriter,
+ ConversionPatternRewriterImpl &rewriterImpl,
+ DenseMap<Value, SmallVector<Value>> &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<Value(Value, Type)> 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<UnrealizedConversionCastOp>();
+ 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<Value, SmallVector<Value>> &inverseMapping) {
+ // As an initial step, compute all of the inserted materializations that we
+ // expect to persist beyond the conversion process.
+ DenseMap<Operation *, UnresolvedMaterializationRewrite *> materializationOps;
+ SetVector<UnresolvedMaterializationRewrite *> 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 75362378daaaaa..156a8a468d5b42 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<tensor = memref<128x64xf16, 3>> to i64
// CHECK-DAG: %[[S1:.+]] = builtin.unrealized_conversion_cast %[[arg1]] : !nvgpu.warpgroup.descriptor<tensor = memref<64x128xf16, 3>> 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]], <m = 64, n = 128, k = 16>, D[<f32>, <one>, <wrapped>], A[<f16>, <one>, <row>], B[<f16>, <one>, <row>] : !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]], <m = 64, n = 128, k = 16>, D[<f32>, <one>, <wrapped>], A[<f16>, <one>, <row>], B[<f16>, <one>, <row>] : !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 ab18ce05e355d3..a192434c5accf8 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<?xf32>) -> memref<?xf32, stri
%0 = bufferization.to_tensor %m : memref<?xf32>
// expected-error @+1 {{failed to legalize unresolved materialization from ('memref<?xf32>') to 'memref<?xf32, strided<[1], offset: ?>>' that remained live after conversion}}
%1 = bufferization.to_memref %0 : memref<?xf32, strided<[1], offset: ?>>
- // expected-note @below{{see existing live user here}}
return %1 : memref<?xf32, strided<[1], offset: ?>>
}
diff --git a/mlir/test/Transforms/test-legalize-type-conversion.mlir b/mlir/test/Transforms/test-legalize-type-conversion.mlir
index f130adff42f8cd..cf2c9f6a8ec441 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 at below {{failed to legalize unresolved materialization from () to 'i16' that remained live after conversion}}
%arg0: i16) {
- // expected-note at 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 at below {{failed to legalize unresolved materialization from ('f64') to 'f16' that remained live after conversion}}
%result = "test.type_producer"() : () -> f16
- // expected-note at 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 at below {{failed to legalize unresolved materialization from ('f64') to 'f16' that remained live after conversion}}
%result = "test.type_producer"() : () -> f16
- // expected-note at 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 at below {{failed to legalize unresolved materialization from ('f64') to 'f16' that remained live after conversion}}
%result = "test.another_type_producer"() : () -> f16
- // expected-note at 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 at below {{failed to legalize unresolved materialization from ('f64') to 'f32' that remained live after conversion}}
+ // expected-note at below {{see existing live user here}}
^bb0(%arg0: f32):
"test.type_consumer"(%arg0) : (f32) -> ()
- // expected-note at below{{see existing live user here}}
"test.return"(%arg0) : (f32) -> ()
}) : () -> ()
return
More information about the Mlir-commits
mailing list