[mlir][transform] Emit silenceable failure for applyToOne

[guraypp]

Summary

This revision emits silenceable failure for applyToOne.

[ftynse]

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

129

Why this change? Diagnostics are rather heavy, we should move them rather than copy.

[guraypp]

Error message does not get printed without this change.
I might doing something wrong.

[ftynse]

There appears to be a bug somewhere in propagation, so let me look into that separately. Please hold this until that is fixed.

[guraypp]

Thanks for looking into that!
This plan sounds good to me. This revision was just to demonstrate the problem.

[ftynse]

https://reviews.llvm.org/D134948 should be the proper fix.

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

1350

Nit: error messages should not start with a capital letter.

mlir/test/Dialect/Linalg/transform-gpu.mlir

188

Nit: @below is more robust than @+1

[guraypp]

I abandon this revision since @ftynse has better solution.

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

LogicalResult silence() {
     if (!diagnostics.empty()) {
       diagnostics.clear();
       result = ::mlir::success();
     }
     return result;
   }
 
   /// Take the diagnostic and silence.
-  SmallVector<Diagnostic> &&takeDiagnostics() {
+  SmallVector<Diagnostic> takeDiagnostics() {
     assert(!diagnostics.empty() && "expected a diagnostic to be present");
     auto guard = llvm::make_scope_exit([&]() { diagnostics.clear(); });
     return std::move(diagnostics);
   }
 
   /// Streams the given values into the last diagnotic.
   /// Expects this object to be a silenceable failure.
   template <typename T>

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

 
 DiagnosedSilenceableFailure
 transform::MapNestedForeachThreadToGpuThreads::applyToOne(
     Operation *target, SmallVectorImpl<Operation *> &results,
     transform::TransformState &state) {
 
   gpu::LaunchOp gpuLaunch = dyn_cast<gpu::LaunchOp>(target);
   if (!gpuLaunch) {
-    target->emitError("Given target is not gpu.launch");
-    return DiagnosedSilenceableFailure::definiteFailure();
+    results.assign({target});
+    return emitSilenceableError() << "Given target is not gpu.launch";
   }
 
   SmallVector<int64_t> blockDim = extractFromI64ArrayAttr(getBlockDim());
   blockDim.resize(/*size=*/3, /*value=*/1);
   SimpleRewriter rewriter(getContext());
   rewriter.setInsertionPoint(target);
   auto walkResult = mlir::linalg::rewriteMapNestedForeachThreadToGpuThreads(
       rewriter, target, blockDim, getSyncAfterDistribute());

mlir/test/Dialect/Linalg/transform-gpu.mlir

 transform.with_pdl_patterns {
 ^bb0(%arg0: !pdl.operation):
   transform.sequence %arg0 failures(propagate) {
   ^bb1(%arg1: !pdl.operation):
     %funcop = transform.structured.match ops{["gpu.launch"]} in %arg0
     transform.structured.map_nested_foreach_thread_to_gpu_threads %funcop { blockDim = [12, 9, 1], syncAfterDistribute = false }
   }
 }
+
+// -----
+
+func.func @foo() -> () {  
+  %1 = linalg.init_tensor [4] : tensor<4xf32>
+  return 
+}
+
+transform.sequence failures(propagate) {
+^bb0(%arg0: !pdl.operation):
+  %funcop = transform.structured.match ops{["linalg.init_tensor"]} in %arg0  
+  // expected-error @+1 {{Given target is not gpu.launch}}
+  %1 = transform.structured.map_nested_foreach_thread_to_gpu_threads %funcop
+}