[mlir] simpler transform dialect silenceable failures

[ftynse]

Summary

Simplify the handling of silenceable failures in the transform dialect.
Previously, the logic of TransformEachOpTrait required that
applyToEach returned a list of null pointers when a silenceable
failure was emitted. This was not done consistently and also crept into
ops without this trait although they did not require it. Handle this
case earlier in the interpreter and homogeneously associated preivously
unset transform dialect values (both handles and parameters) with empty
lists of the matching kind. Ignore the results of applyToEach for the
targets for which it produced a silenceable failure. As a result, one
never needs to set results to lists containing nulls. Furthermore, the
objects associated with transform dialect values must never be null.

Depends On D140980

[nicolasvasilache]

Great, thanks!

mlir/include/mlir/Dialect/Transform/IR/TransformInterfaces.h

private:
   /// be associated with parameters.
   ArrayRef<TransformState::Param> getParams(unsigned resultNumber) const;
 
   /// Returns `true` if the result identified by its number in the list of
   /// operation results is associated with a list of parameters, `false` if it
   /// is associated with the list of payload IR operations.
   bool isParam(unsigned resultNumber) const;
 
+  /// Returns `true` if the result identified by its number in the list of
+  /// operation results is associated with something.
+  bool isSet(unsigned resultNumber) const;
+
   /// Storage for pointers to payload IR ops that are associated with results of
   /// a transform IR op. `segments` contains as many entries as the transform IR
   /// op has results, even if some of them are not associated with payload IR
   /// operations. Each entry is a reference to a contiguous segment in the
   /// `operations` list that contains the pointers to operations. This allows
   /// for operations to be stored contiguously without nested vectors and for
   /// different segments to be set in any order.
   SmallVector<ArrayRef<Operation *>, 2> segments;
if (!specificOp) {
       continue;
     }
 
     ApplyToEachResultList partialResults;
     partialResults.reserve(expectedNumResults);
     Location specificOpLoc = specificOp->getLoc();
     DiagnosedSilenceableFailure res =
         transformOp.applyToOne(specificOp, partialResults, state);
-    if (res.isDefiniteFailure() ||
-        failed(detail::checkApplyToOne(transformOp, specificOpLoc,
-                                       partialResults))) {
+    if (res.isDefiniteFailure())
       return DiagnosedSilenceableFailure::definiteFailure();
-    }
 
-    if (res.isSilenceableFailure())
+    if (res.isSilenceableFailure()) {
       res.takeDiagnostics(silenceableStack);
-    else
-      results.push_back(std::move(partialResults));
+      continue;
+    }
+
+    if (failed(detail::checkApplyToOne(transformOp, specificOpLoc,
+                                       partialResults))) {
+      return DiagnosedSilenceableFailure::definiteFailure();
+    }
   }
   if (!silenceableStack.empty()) {
     return DiagnosedSilenceableFailure::silenceableFailure(
         std::move(silenceableStack));
   }
   return DiagnosedSilenceableFailure::success();
 }
 

mlir/lib/Dialect/GPU/TransformOps/GPUTransformOps.cpp

 transform::MapForeachToBlocks::applyToOne(Operation *target,
                                           ApplyToEachResultList &results,
                                           transform::TransformState &state) {
   LaunchOp gpuLaunch = dyn_cast<LaunchOp>(target);
   TrivialPatternRewriter rewriter(getContext());
   auto transformOp = cast<TransformOpInterface>(getOperation());
 
   if (!getGenerateGpuLaunch() && !gpuLaunch) {
-    results.push_back(target);
     DiagnosedSilenceableFailure diag =
         emitSilenceableError()
         << "Given target is not gpu.launch, set `generate_gpu_launch` "
            "attribute";
     diag.attachNote(target->getLoc()) << "when applied to this payload op";
     return diag;
   }
 
   scf::ForeachThreadOp topLevelForeachThreadOp;
   DiagnosedSilenceableFailure diag =
       mlir::transform::gpu::findTopLevelForeachThreadOp(
           target, topLevelForeachThreadOp, transformOp);
   if (!diag.succeeded()) {
-    results.push_back(target);
     diag.attachNote(target->getLoc()) << "when applied to this payload op";
     return diag;
   }
 
   OpBuilder::InsertionGuard guard(rewriter);
   rewriter.setInsertionPoint(topLevelForeachThreadOp);
 
   // Generate gpu launch here and move the foreach_thread inside
   if (getGenerateGpuLaunch()) {
     DiagnosedSilenceableFailure diag =
         createGpuLaunch(rewriter, target->getLoc(), transformOp, gpuLaunch);
     if (!diag.succeeded()) {
-      results.push_back(target);
       return diag;
     }
     rewriter.setInsertionPointToStart(&gpuLaunch.getBody().front());
     Operation *newForeachThreadOp = rewriter.clone(*topLevelForeachThreadOp);
     rewriter.eraseOp(topLevelForeachThreadOp);
     topLevelForeachThreadOp = cast<scf::ForeachThreadOp>(newForeachThreadOp);
   }
 

 
 //===----------------------------------------------------------------------===//
 // MapNestedForeachToThreads
 //===----------------------------------------------------------------------===//
 
 /// Searches `scf.foreach_thread` ops nested under `target` and maps each such
 /// op to GPU threads. Mapping is one-to-one and the induction variables of
 /// `scf.foreach_thread` are rewritten to gpu.thread_id according to the
-/// thread_dim_apping attribute. Sibling `scf.foreach_thread` are supported in
+/// thread_dim_mapping attribute. Sibling `scf.foreach_thread` are supported in
 /// which case, the union of the number of threads is computed and may result
 /// in predication. Dynamic, `scf.foreach_thread` trip counts are currently
 /// not supported. Dynamic block dim sizes are currently not supported.
 static DiagnosedSilenceableFailure rewriteOneForeachThreadToGpuThreads(
     RewriterBase &rewriter, scf::ForeachThreadOp foreachThreadOp,
     const SmallVectorImpl<int64_t> &globalBlockDims, bool syncAfterDistribute,
     std::optional<TransformOpInterface> transformOp,
     const ArrayRef<DeviceMappingAttrInterface> &threadMappingAttributes) {

 }
 
 DiagnosedSilenceableFailure transform::MapNestedForeachToThreads::applyToOne(
     Operation *target, ApplyToEachResultList &results, TransformState &state) {
   LaunchOp gpuLaunch = dyn_cast<LaunchOp>(target);
   auto transformOp = cast<TransformOpInterface>(getOperation());
 
   if (!gpuLaunch) {
-    results.push_back(target);
     return emitSilenceableError() << "Given target is not gpu.launch";
   }
 
   SmallVector<int64_t> blockDim = extractFromI64ArrayAttr(getBlockDim());
   blockDim.resize(/*size=*/3, /*value=*/1);
 
   DiagnosedSilenceableFailure diag =
       checkGpuLimits(transformOp, std::nullopt, std::nullopt, std::nullopt,
                      blockDim[0], blockDim[1], blockDim[2]);
   if (diag.isSilenceableFailure()) {
-    results.push_back(target);
     diag.attachNote(getLoc()) << getBlockDimAttrName() << " is very large";
     return diag;
   }
 
   MLIRContext *ctx = getContext();
   TrivialPatternRewriter rewriter(ctx);
   rewriter.setInsertionPoint(target);
 

mlir/lib/Dialect/Linalg/TransformOps/LinalgTransformOps.cpp

#define DOWNSCALE_NORMAL(a, b) DOWNSCALE(DOWNSCALE_CALL(a, b))
   DOWNSCALE_NORMAL(PoolingNhwcMaxUnsignedOp, PoolingNwcMaxUnsignedOp)
   DOWNSCALE_NORMAL(PoolingNhwcMinOp, PoolingNwcMinOp)
   DOWNSCALE_NORMAL(PoolingNhwcMinUnsignedOp, PoolingNwcMinUnsignedOp)
   DOWNSCALE_NORMAL(PoolingNchwMaxOp, PoolingNcwMaxOp)
   DOWNSCALE(DownscaleDepthwiseConv2DNhwcHwcOp)
 #undef DOWNSCALE_NORMAL
 #undef DOWNSCALE_CALL
 #undef DOWNSCALE
-  results.assign(1, nullptr);
   return emitDefaultSilenceableFailure(target);
 }
 //===----------------------------------------------------------------------===//
 // FuseOp
 //===----------------------------------------------------------------------===//
 
 /// Apply a tiling transformation to all payload ops and store both the
 /// tiled operation as well as the created tile loops.
auto getNextProducer = [&]() -> FailureOr<Operation *> {
     }
     return failure();
   };
 
   IRRewriter rewriter(getContext());
   while (!remainingProducers.empty()) {
     auto nextProducer = getNextProducer();
     if (failed(nextProducer)) {
-      results.set(getFusedOp().cast<OpResult>(), ArrayRef<Operation *>());
       Diagnostic diag(containingOp->getLoc(), DiagnosticSeverity::Remark);
       diag << "could not find next producer to fuse into container";
       return DiagnosedSilenceableFailure::silenceableFailure(std::move(diag));
     }
 
     Operation *producerOp = *nextProducer;
 
     // Default diagnostic, to be complemented with more failure information.
while (!remainingProducers.empty()) {
     Operation *cloned =
         cloneAndFuseFirstUse(rewriter, diag, producerOp, containingOp);
     if (cloned) {
       LLVM_DEBUG(llvm::dbgs() << "\nFused an use by cloning\n"
                               << *containingOp);
       fusedOps.push_back(cloned);
       continue;
     }
-    results.set(getFusedOp().cast<OpResult>(), ArrayRef<Operation *>());
     return DiagnosedSilenceableFailure::silenceableFailure(std::move(diag));
   }
 
   results.set(getFusedOp().cast<OpResult>(), fusedOps);
   return DiagnosedSilenceableFailure::success();
 }
 
 //===----------------------------------------------------------------------===//
if (isa<GenericOp>(target)) {
     results.push_back(target);
     return DiagnosedSilenceableFailure::success();
   }
   FailureOr<LinalgOp> generic = tryApply<LinalgGeneralizationPattern>(target);
   if (succeeded(generic)) {
     results.push_back(generic->getOperation());
     return DiagnosedSilenceableFailure::success();
   }
-  results.assign(1, nullptr);
   return emitDefaultSilenceableFailure(target);
 }
 
 //===----------------------------------------------------------------------===//
 // InterchangeOp
 //===----------------------------------------------------------------------===//
 
 DiagnosedSilenceableFailure
transform::MatchOp::apply(transform::TransformResults &results,
                           transform::TransformState &state) {
   llvm::StringSet<> strs;
   if (getOps().has_value())
     strs.insert(getOps()->getAsValueRange<StringAttr>().begin(),
                 getOps()->getAsValueRange<StringAttr>().end());
 
   ArrayRef<Operation *> payloadOps = state.getPayloadOps(getTarget());
   if (payloadOps.size() != 1) {
-    results.set(getResult().cast<OpResult>(), {});
     return emitDefiniteFailure("requires exactly one target handle");
   }
 
   SmallVector<Operation *> res;
   auto matchFun = [&](Operation *op) {
     if (getOps().has_value() && !strs.contains(op->getName().getStringRef()))
       return;
 
static ParseResult parseMultitileSizesTypes(OpAsmParser &parser,
   return success();
 }
 
 DiagnosedSilenceableFailure transform::MultiTileSizesOp::applyToOne(
     LinalgOp target, transform::ApplyToEachResultList &results,
     TransformState &state) {
   if (getLowSize().getType().isa<TransformParamTypeInterface>()) {
     if (target.hasDynamicShape()) {
-      results.assign(
-          ArrayRef<Attribute>({Attribute(), Attribute(), Attribute()}));
       auto diag = emitSilenceableError()
                   << "cannot compute parametric tile sizes for dynamically "
                      "shaped payload op";
       diag.attachNote(target->getLoc()) << "payload op";
       return diag;
     }
 
     FailureOr<StaticMultiSizeSpecification> spec = computeStaticMultiTileSizes(
         target, getDimension(), getTargetSize(), getDivisor());
     if (failed(spec)) {
-      results.assign(
-          ArrayRef<Attribute>({Attribute(), Attribute(), Attribute()}));
       return emitSilenceableError()
              << "failed to compute multi-size tiling sizes";
     }
 
     Builder builder(target.getContext());
     results.assign(llvm::map_range(
         ArrayRef<int64_t>({spec->lowTileSize, spec->highTileSize,
                            spec->lowTileSize * spec->lowTripCount}),
transform::PadOp::applyToOne(linalg::LinalgOp target,
 
   FailureOr<LinalgOp> result =
       tryApply<LinalgPaddingPattern>(target, paddingOptions);
   if (succeeded(result)) {
     results.push_back(result->getOperation());
     return DiagnosedSilenceableFailure::success();
   }
 
-  results.assign(1, nullptr);
   return emitDefaultSilenceableFailure(target);
 }
 
 LogicalResult transform::PadOp::verify() {
   SmallVector<int64_t> packPaddings =
       extractFromI64ArrayAttr(getPackPaddings());
   if (any_of(packPaddings, [](int64_t packPadding) {
         return packPadding != 0 && packPadding != 1;
if (getDynamicSplitPoint().getType().isa<TransformHandleTypeInterface>()) {
             }
             return OpFoldResult(op->getResult(0));
           }));
     } else {
       splitPoints = llvm::to_vector(
           llvm::map_range(state.getParams(getDynamicSplitPoint()),
                           [](Attribute attr) { return OpFoldResult(attr); }));
     }
-    if (diag.isSilenceableFailure()) {
-      results.set(getFirst().cast<OpResult>(), {});
-      results.set(getSecond().cast<OpResult>(), {});
+    if (diag.isSilenceableFailure())
       return diag;
-    }
 
     if (splitPoints.size() != payload.size()) {
       return emitDefiniteFailure()
              << "expected the dynamic split point handle to point to as "
                 "many operations ("
              << splitPoints.size() << ") as the target handle ("
              << payload.size() << ")";
     }
   } else {
     splitPoints.resize(payload.size(),
                        rewriter.getIndexAttr(getStaticSplitPoint()));
   }
 
   // Split each target operation.
   SmallVector<Operation *> first, second;
   Operation *noSecondPart = nullptr;
   for (const auto &pair : llvm::zip(payload, splitPoints)) {
     Operation *target = std::get<0>(pair);
     auto linalgOp = dyn_cast<LinalgOp>(target);
     if (!linalgOp) {
       auto diag = emitSilenceableError() << "only applies to structured ops";
       diag.attachNote(target->getLoc()) << "target op";
-      results.set(getFirst().cast<OpResult>(), {});
-      results.set(getSecond().cast<OpResult>(), {});
       return diag;
     }
 
     if (getDimension() >= linalgOp.getNumLoops()) {
       auto diag = emitSilenceableError() << "dimension " << getDimension()
                                          << " does not exist in target op";
       diag.attachNote(target->getLoc()) << "target op";
-      results.set(getFirst().cast<OpResult>(), {});
-      results.set(getSecond().cast<OpResult>(), {});
       return diag;
     }
 
     rewriter.setInsertionPoint(linalgOp);
     std::tie(first.emplace_back(), second.emplace_back()) = linalg::splitOp(
         rewriter, cast<TilingInterface>(linalgOp.getOperation()),
         getDimension(), std::get<1>(pair));
 
     // Propagate errors.
     if (!first.back() && !second.back()) {
       auto diag = emitDefiniteFailure() << "internal failure in splitting";
       diag.attachNote(target->getLoc()) << "target op";
       return diag;
     }
 
     // Do not add null second parts.
     if (!second.back()) {
       noSecondPart = target;
       second.pop_back();
     }
   }
 
   if (second.size() != first.size() && !second.empty()) {
-    results.set(getFirst().cast<OpResult>(), {});
-    results.set(getSecond().cast<OpResult>(), {});
     auto diag =
         emitSilenceableError()
         << "splitting does not produce the second part for a subset of targets";
     diag.attachNote() << "expected splitting to produce the second part of all "
                          "or none of the targets";
     diag.attachNote(noSecondPart->getLoc())
         << "first target with no second part";
     return diag;
transform::TileReductionUsingForeachThreadOp::applyToOne(
   SmallVector<OpFoldResult> tileSizes =
       getAsOpFoldResult(rewriter.getI64ArrayAttr(getTileSizes()));
   FailureOr<linalg::ForeachThreadReductionTilingResult> result =
       linalg::tileReductionUsingForeachThread(
           rewriter, cast<PartialReductionOpInterface>(target.getOperation()),
           numThreads, tileSizes, getMapping());
 
   if (failed(result)) {
-    results.assign(4, nullptr);
     auto diag = emitSilenceableError() << "could not tile reduction";
     diag.attachNote(target.getLoc()) << "target operation";
     return diag;
   }
   results.push_back(result->loops);
   results.push_back(result->initialOp);
   results.push_back(result->parallelTiledOp);
   results.push_back(result->mergeOp);
if (transformValue.getType().isa<ParamType>()) {
       dynamicSizeProducers.push_back({});
       ArrayRef<Attribute> params = state.getParams(transformValue);
       paramSizes.push_back(
           llvm::to_vector(llvm::map_range(params, [](Attribute attr) {
             return attr.cast<IntegerAttr>().getValue().getSExtValue();
           })));
 
       if (paramSizes.back().size() != targets.size()) {
-        for (OpResult r : getResults())
-          transformResults.set(r, {});
         DiagnosedSilenceableFailure diag =
             emitSilenceableError()
             << "expected as many parameter values ("
             << dynamicSizeProducers.back().size() << ") as target ops ("
             << targets.size() << ")";
         diag.attachNote(transformValue.getLoc()) << "for this parameter";
         return diag;
       }
 
       continue;
     }
     paramSizes.push_back({});
     dynamicSizeProducers.push_back(state.getPayloadOps(transformValue));
 
     if (dynamicSizeProducers.back().size() != targets.size()) {
-      for (OpResult r : getResults())
-        transformResults.set(r, {});
       DiagnosedSilenceableFailure diag =
           emitSilenceableError()
           << "expected as many dynamic size-producing operations ("
           << dynamicSizeProducers.back().size() << ") as target ops ("
           << targets.size() << ")";
       diag.attachNote(transformValue.getLoc()) << "for this handle";
       return diag;
     }
 
     for (Operation *op : dynamicSizeProducers.back()) {
       if (op->getNumResults() == 1 &&
           op->getResult(0).getType().isa<IndexType>())
         continue;
 
-      for (OpResult r : getResults())
-        transformResults.set(r, {});
       DiagnosedSilenceableFailure diag =
           emitSilenceableError() << "expected sizes to be produced by ops "
                                     "with a single index-type result";
       diag.attachNote(op->getLoc()) << "size producer op";
       diag.attachNote(transformValue.getLoc()) << "for this handle";
       return diag;
     }
   }
status = getPackedTileSizes()
                      state, transformOp, mixedTileSizes, getMixedTileSizes());
   if (!status.succeeded())
     return status;
 
   DiagnosedSilenceableFailure diag = tileToForeachThreadOpImpl(
       rewriter, state, transformOp, targets, mixedNumThreads, mixedTileSizes,
       getMapping(), tileOps, tiledOps);
 
-  if (!diag.succeeded()) {
-    transformResults.set(getForeachThreadOp().cast<OpResult>(), {});
-    transformResults.set(getTiledOp().cast<OpResult>(), {});
+  if (!diag.succeeded())
     return diag;
-  }
 
   transformResults.set(getForeachThreadOp().cast<OpResult>(), tileOps);
   transformResults.set(getTiledOp().cast<OpResult>(), tiledOps);
 
   return DiagnosedSilenceableFailure::success();
 }
 
 void transform::TileToForeachThreadOp::getEffects(

mlir/lib/Dialect/SCF/TransformOps/SCFTransformOps.cpp

for (unsigned i = 0, e = getNumLoops(); i < e; ++i) {
       if (!loop) {
         DiagnosedSilenceableFailure diag =
             emitSilenceableError()
             << "could not find an '"
             << (getAffine() ? AffineForOp::getOperationName()
                             : scf::ForOp::getOperationName())
             << "' parent";
         diag.attachNote(target->getLoc()) << "target op";
-        results.set(getResult().cast<OpResult>(), {});
         return diag;
       }
       current = loop;
     }
     parents.insert(loop);
   }
   results.set(getResult().cast<OpResult>(), parents.getArrayRef());
   return DiagnosedSilenceableFailure::success();
for (Operation *target : state.getPayloadOps(getTarget())) {
     Location location = target->getLoc();
     Operation *symbolTableOp = SymbolTable::getNearestSymbolTable(target);
     TrivialPatternRewriter rewriter(getContext());
     scf::ExecuteRegionOp exec = wrapInExecuteRegion(rewriter, target);
     if (!exec) {
       DiagnosedSilenceableFailure diag = emitSilenceableError()
                                          << "failed to outline";
       diag.attachNote(target->getLoc()) << "target op";
-      results.set(getTransformed().cast<OpResult>(), {});
       return diag;
     }
     func::CallOp call;
     FailureOr<func::FuncOp> outlined = outlineSingleBlockRegion(
         rewriter, location, exec.getRegion(), getFuncName(), &call);
 
     if (failed(outlined))
       return emitDefaultDefiniteFailure(target);
transform::LoopPipelineOp::applyToOne(scf::ForOp target,
   TrivialPatternRewriter rewriter(getContext());
   rewriter.setInsertionPoint(target);
   FailureOr<scf::ForOp> patternResult =
       pattern.returningMatchAndRewrite(target, rewriter);
   if (succeeded(patternResult)) {
     results.push_back(*patternResult);
     return DiagnosedSilenceableFailure::success();
   }
-  results.assign(1, nullptr);
   return emitDefaultSilenceableFailure(target);
 }
 
 //===----------------------------------------------------------------------===//
 // LoopUnrollOp
 //===----------------------------------------------------------------------===//
 
 DiagnosedSilenceableFailure

mlir/lib/Dialect/Transform/IR/TransformInterfaces.cpp

transform::TransformState::applyTransform(TransformOpInterface transform) {
   transform::TransformResults results(transform->getNumResults());
   // Compute the result but do not short-circuit the silenceable failure case as
   // we still want the handles to propagate properly so the "suppress" mode can
   // proceed on a best effort basis.
   DiagnosedSilenceableFailure result(transform.apply(results, *this));
   if (result.isDefiniteFailure())
     return result;
 
+  // If a silenceable failure was produced, some results may be unset, set them
+  // to empty lists.
+  if (result.isSilenceableFailure()) {
+    for (OpResult opResult : transform->getResults()) {
+      if (results.isSet(opResult.getResultNumber()))
+        continue;
+
+      if (opResult.getType().isa<TransformParamTypeInterface>())
+        results.setParams(opResult, {});
+      else
+        results.set(opResult, {});
+    }
+  }
+
   // Remove the mapping for the operand if it is consumed by the operation. This
   // allows us to catch use-after-free with assertions later on.
   auto memEffectInterface =
       cast<MemoryEffectOpInterface>(transform.getOperation());
   SmallVector<MemoryEffects::EffectInstance, 2> effects;
   for (OpOperand &target : transform->getOpOperands()) {
     effects.clear();
     memEffectInterface.getEffectsOnValue(target.get(), effects);

 }
 
 bool transform::TransformResults::isParam(unsigned resultNumber) const {
   assert(resultNumber < paramSegments.size() &&
          "querying association for a non-existent handle");
   return paramSegments[resultNumber].data() != nullptr;
 }
 
+bool transform::TransformResults::isSet(unsigned resultNumber) const {
+  assert(resultNumber < paramSegments.size() &&
+         "querying association for a non-existent handle");
+  return paramSegments[resultNumber].data() != nullptr ||
+         segments[resultNumber].data() != nullptr;
+}
+
 //===----------------------------------------------------------------------===//
 // Utilities for TransformEachOpTrait.
 //===----------------------------------------------------------------------===//
 
 LogicalResult
 transform::detail::checkApplyToOne(Operation *transformOp,
                                    Location payloadOpLoc,
                                    const ApplyToEachResultList &partialResult) {
   Location transformOpLoc = transformOp->getLoc();
   StringRef transformOpName = transformOp->getName().getStringRef();
   unsigned expectedNumResults = transformOp->getNumResults();
-  // TODO: encode this implicit must always produce `expectedNumResults`
-  // and nullptr is fine with a proper trait.
+
+  // Reuse the emission of the diagnostic note.
+  auto emitDiag = [&]() {
+    auto diag = mlir::emitError(transformOpLoc);
+    diag.attachNote(payloadOpLoc) << "when applied to this op";
+    return diag;
+  };
+
   if (partialResult.size() != expectedNumResults) {
-    auto diag = mlir::emitError(transformOpLoc, "applications of ")
-                << transformOpName << " expected to produce "
-                << expectedNumResults << " results (actually produced "
+    auto diag = emitDiag() << "application of " << transformOpName
+                           << " expected to produce " << expectedNumResults
+                           << " results (actually produced "
                            << partialResult.size() << ").";
     diag.attachNote(transformOpLoc)
-        << "If you need variadic results, consider a generic `apply` "
+        << "if you need variadic results, consider a generic `apply` "
         << "instead of the specialized `applyToOne`.";
-    diag.attachNote(transformOpLoc)
-        << "Producing " << expectedNumResults << " null results is "
-        << "allowed if the use case warrants it.";
-    diag.attachNote(payloadOpLoc) << "when applied to this op";
-    return failure();
-  }
-
-  // Check that all is null or none is null
-  // TODO: relax this behavior and encode with a proper trait.
-  if (llvm::any_of(
-          partialResult,
-          [](llvm::PointerUnion<Operation *, Attribute> ptr) { return ptr; }) &&
-      llvm::any_of(partialResult,
-                   [](llvm::PointerUnion<Operation *, Attribute> ptr) {
-                     return !ptr;
-                   })) {
-    auto diag = mlir::emitError(transformOpLoc, "unexpected application of ")
-                << transformOpName
-                << " produces both null and non null results.";
-    diag.attachNote(payloadOpLoc) << "when applied to this op";
     return failure();
   }
 
   // Check that the right kind of value was produced.
   for (const auto &[ptr, res] :
        llvm::zip(partialResult, transformOp->getResults())) {
+    if (ptr.isNull()) {
+      return emitDiag() << "null result #" << res.getResultNumber()
+                        << " produced";
+    }
     if (ptr.is<Operation *>() &&
         !res.getType().template isa<TransformHandleTypeInterface>()) {
-      mlir::emitError(transformOpLoc)
-          << "applications of " << transformOpName
+      return emitDiag() << "application of " << transformOpName
                         << " expected to produce an Attribute for result #"
                         << res.getResultNumber();
-      return failure();
     }
     if (ptr.is<Attribute>() &&
         !res.getType().template isa<TransformParamTypeInterface>()) {
-      mlir::emitError(transformOpLoc)
-          << "applications of " << transformOpName
+      return emitDiag() << "application of " << transformOpName
                         << " expected to produce an Operation * for result #"
                         << res.getResultNumber();
-      return failure();
     }
   }
   return success();
 }
 
 void transform::detail::setApplyToOneResults(
     Operation *transformOp, TransformResults &transformResults,
     ArrayRef<ApplyToEachResultList> results) {

mlir/lib/Dialect/Transform/IR/TransformOps.cpp

Operation *producer =
             ? nullptr
             : target->getOperand(operandNumber).getDefiningOp();
     if (!producer) {
       DiagnosedSilenceableFailure diag =
           emitSilenceableError()
           << "could not find a producer for operand number: " << operandNumber
           << " of " << *target;
       diag.attachNote(target->getLoc()) << "target op";
-      results.set(getResult().cast<OpResult>(),
-                  SmallVector<mlir::Operation *>{});
       return diag;
     }
     producers.push_back(producer);
   }
   results.set(getResult().cast<OpResult>(), producers);
   return DiagnosedSilenceableFailure::success();
 }
 

 
 DiagnosedSilenceableFailure
 transform::SplitHandlesOp::apply(transform::TransformResults &results,
                                  transform::TransformState &state) {
   int64_t numResultHandles =
       getHandle() ? state.getPayloadOps(getHandle()).size() : 0;
   int64_t expectedNumResultHandles = getNumResultHandles();
   if (numResultHandles != expectedNumResultHandles) {
-    // Failing case needs to propagate gracefully for both suppress and
-    // propagate modes.
-    for (int64_t idx = 0; idx < expectedNumResultHandles; ++idx)
-      results.set(getResults()[idx].cast<OpResult>(), {});
     // Empty input handle corner case: always propagates empty handles in both
     // suppress and propagate modes.
     if (numResultHandles == 0)
       return DiagnosedSilenceableFailure::success();
     // If the input handle was not empty and the number of result handles does
     // not match, this is a legit silenceable error.
     return emitSilenceableError()
            << getHandle() << " expected to contain " << expectedNumResultHandles

mlir/test/Dialect/Transform/test-interpreter.mlir

pdl.pattern @some : benefit(1) {
     %1 = pdl.types
     %2 = pdl.operation "op"(%0 : !pdl.range<value>) -> (%1 : !pdl.range<type>)
     pdl.rewrite %2 with "transform.dialect"
   }
 
   transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb0(%arg1: !pdl.operation):
     %0 = pdl_match @some in %arg1 : (!pdl.operation) -> !pdl.operation
-    // expected-error @below {{applications of transform.test_wrong_number_of_results expected to produce 3 results (actually produced 1).}}
-    // expected-note @below {{If you need variadic results, consider a generic `apply` instead of the specialized `applyToOne`.}}
-    // expected-note @below {{Producing 3 null results is allowed if the use case warrants it.}}
+    // expected-error @below {{application of transform.test_wrong_number_of_results expected to produce 3 results (actually produced 1).}}
+    // expected-note @below {{if you need variadic results, consider a generic `apply` instead of the specialized `applyToOne`.}}
     transform.test_wrong_number_of_results %0
   }
 }
 
 // -----
 
 func.func @foo() {
   "op" () : () -> ()
pdl.pattern @some : benefit(1) {
     %1 = pdl.types
     %2 = pdl.operation "op"(%0 : !pdl.range<value>) -> (%1 : !pdl.range<type>)
     pdl.rewrite %2 with "transform.dialect"
   }
 
   transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb0(%arg1: !pdl.operation):
     %0 = pdl_match @some in %arg1 : (!pdl.operation) -> !pdl.operation
-    // expected-error @below {{applications of transform.test_wrong_number_of_multi_results expected to produce 1 results (actually produced 0)}}
-    // expected-note @below {{If you need variadic results, consider a generic `apply` instead of the specialized `applyToOne`.}}
-    // expected-note @below {{Producing 1 null results is allowed if the use case warrants it.}}
+    // expected-error @below {{application of transform.test_wrong_number_of_multi_results expected to produce 1 results (actually produced 0)}}
+    // expected-note @below {{if you need variadic results, consider a generic `apply` instead of the specialized `applyToOne`.}}
     transform.test_wrong_number_of_multi_results %0
   }
 }
 
 // -----
 
 func.func @foo() {
   "op" () : () -> ()
pdl.pattern @some : benefit(1) {
     %1 = pdl.types
     %2 = pdl.operation "op"(%0 : !pdl.range<value>) -> (%1 : !pdl.range<type>)
     pdl.rewrite %2 with "transform.dialect"
   }
 
   transform.sequence %arg0 : !pdl.operation failures(propagate) {
   ^bb0(%arg1: !pdl.operation):
     %0 = pdl_match @some in %arg1 : (!pdl.operation) -> !pdl.operation
-    // expected-error @below {{unexpected application of transform.test_mixed_null_and_non_null_results produces both null and non null results.}}
+    // expected-error @below {{null result #0 produced}}
     transform.test_mixed_null_and_non_null_results %0
   }
 }
 
 // -----
 
 // Expecting to match all operations by merging the handles that matched addi
 // and subi separately.
func.func private @three_test_ops(%arg0: i32) {
   "test.op_a"(%arg0) { attr = 0 : i32} : (i32) -> i32
   "test.op_a"(%arg0) { attr = 0 : i32} : (i32) -> i32
   "test.op_a"(%arg0) { attr = 0 : i32} : (i32) -> i32
   return
 }
 
 // -----
 
+// expected-note @below {{when applied to this op}}
+module {
   transform.sequence failures(propagate) {
   ^bb0(%arg0: !transform.any_op):
     // expected-error @below {{expected to produce an Operation * for result #0}}
     transform.test_produce_transform_param_or_forward_operand %arg0
       { first_result_is_param }
       : (!transform.any_op) -> (!transform.any_op, !transform.param<i64>)
   }
+}
 
 // -----
 
 // expected-note @below {{when applied to this op}}
 module {
   transform.sequence failures(propagate) {
   ^bb0(%arg0: !transform.any_op):
-    // expected-error @below {{produces both null and non null results}}
+    // expected-error @below {{null result #0 produced}}
     transform.test_produce_transform_param_or_forward_operand %arg0
       { first_result_is_null }
       : (!transform.any_op) -> (!transform.any_op, !transform.param<i64>)
   }
 }
 
 // -----
 
+// expected-note @below {{when applied to this op}}
+module {
   transform.sequence failures(propagate) {
   ^bb0(%arg0: !transform.any_op):
     // expected-error @below {{expected to produce an Attribute for result #1}}
     transform.test_produce_transform_param_or_forward_operand %arg0
       { second_result_is_handle }
       : (!transform.any_op) -> (!transform.any_op, !transform.param<i64>)
   }
+}
 
 // -----
 
 transform.sequence failures(propagate) {
 ^bb0(%arg0: !transform.any_op):
   // expected-error @below {{attempting to assign a null payload op to this transform value}}
   %0 = transform.test_produce_null_payload : !transform.any_op
 }
 
 // -----
 
 transform.sequence failures(propagate) {
 ^bb0(%arg0: !transform.any_op):
   // expected-error @below {{attempting to assign a null parameter to this transform value}}
   %0 = transform.test_produce_null_param : !transform.param<i64>
 }