[Cherry-Pick] gfx950 improvements and bug fixes (#846)

* [AMD] Avoid async load to pipeline for less than 32bit load (triton-lang#7250)

We can only use AsyncCopy if the final load width can be >= 4 bytes. 
`triton::canBeConvertedToAsyncLoad` checks that the vecSize of the
source is large enough. Additionally we need to ensure the register to
shared layout (blocked+shared) does have enough contiguous elements
since we cannot scatter into LDS.

Before this PR we will abort compilation instead of falling back to
pipelining through registers.

* [AMD] Pipeline small tensors w/ registers only on GFX950 (triton-lang#7171)

Fixes a perf regression on gfx942 but preserves functionality for
gfx950 (and above).

* Reland "[AMD] Optimize reduction with v_permlane intrinsics in GFX950" (triton-lang#7321)

triton-lang#7291 fixed the
LLVM issue that caused correctness problems. Now
we can reland this patch.

* [Pipeliner] Expose core pipeliner utilities to reuse in AMD backend (triton-lang#7222)

This PR exposes (via header) core pipelining utilities/helpers:

```c++
bool hasGpuBarriers(scf::ForOp forOp);
bool isSafeToPipeline(scf::ForOp forOp);
llvm::MapVector<Operation *, std::pair<int, Operation *>>
loadOpsToIndirectionLevel(scf::ForOp forOp, bool pipelineWithoutDot,
                          triton::ModuleAxisInfoAnalysis &axisInfoAnalysis,
                          int numStages, bool filterSmall = true);
void scheduleDistanceOneDependencies(scf::ForOp forOp,
                                     CoarseSchedule &schedule);
void scheduleRemainingToLastStage(scf::ForOp forOp, CoarseSchedule &schedule,
                                  CoarseSchedule::Cluster afterPrologue);
```

They are directly useable by AMD's pipeliner. 

Note, this is basically NFC for AMD because AMD's pipeliner simply had
copy-paste of the same functions from ~last year.


Small API changes:

1. On NV we do not pipeline small loads (vec width < 32b). On AMD we do.
The choice is made inside `isPipeliningBeneficial` inside
`loadOpsToIndirectionLevel`. To support AMD I have added a flag
`filterSmall`.
2. On AMD the load `use`s (computed as a matter of course in
`loadOpsToIndirectionLevel`) are used (no pun intended) whereas on NV
they are not. To support AMD I keep those `use`s in the
`llvm::MapVector<Operation *, std::pair<int, Operation *>>` return from
`loadOpsToIndirectionLevel`.

These two small changes are the only "non-NFC" changes.

* [AMD] Retire local prefetch schedule hint variant (triton-lang#7395)

This variant was from some prior experiments. We have a better way
to implement later.

* [AMD] Retire TritonAMDGPU_OpIdxAttr and TritonAMDGPU_InstCounter  (triton-lang#7476)

triton-lang#7395 retired the local
prefetch schedule variant. This made `TritonAMDGPU_OpIdxAttr` and
`TritonAMDGPU_InstCounter` unused which are removed by this PR.

* [AMD][NFC] Split createAndSchedule* in stream pipeliner(triton-lang#7514)

Splits `createAndScheduleAsyncCopy` and `createAndScheduleStreamCopy` to
make it reusable if we want to schedule the ops differently in a future
PR.

* [AMD] Refactor StreamPipeliner to use more common functions (triton-lang#7526)

Further refactoring of Streampipeliner.cpp to use more common pipeliner
functionality: `triton::createAllocation`,
`triton::createSingleBufferView`, `triton::replaceWithSharedLoad` and a
bit of general cleanup.

Overall NFC except:
- The order of LocalDealloc is reversed now
- The memdesc of the subview additionally includes the allocSize

Also we had no lit test checking that the LocalLoad consumes the
AsyncToken so I adjusted one to include the check.

* [AMD] NFC: Refactor stream pipeliner to better encapsulate functions (triton-lang#7540)

Mostly moves code around to reduce the dependencies between functions
and further splits up functions doing more than one thing
(`createAndSchedule*`,` preprocessAndBuildSchedule`). This will also
allow us to use more common pipeliner functionality in a future PR, e.g.
`createAsyncCopy`.

* [FA] Set vecSize=nonKDim for V shared layout to avoid bank conflicts

I'll submit a PR upstream later.

* [GEMM] Add combine dot_scaled and addF

* [AMD][NFC] Consolidate initialization in initSchedule for pipeliner (triton-lang#7556)

Moves all initializations of stages to `initSchedule`. Missed this one
in the last PRs.

* [AMD] NFC: Drop version minor for AMD MFMA layout (triton-lang#7285)

AMD's MFMA layout does not need version minor information like NVIDIA.
It always defaults to 0 in the current codebase. The PR drops version
minor and change to a single `version` parameter for MFMA layout.

* [AMD] Add tilesPerWarp parameter to mfma layout (triton-lang#7283)

This PR introduces the tilesPerWarp parameter to the MFMA layout.
Previously, the MFMA layout assumed that each warp within a CTA tile
computed a single MFMA tile.
When the tensor was larger than a single CTA tile, these tiles were
repeated across the tensor.
In this setup, the output tiles computed by each wave were strided by
the number of warps
per CTA in both row and column dimensions.

For instance, with 16 MFMA tiles and warpsPerCTA = [2, 2], the
distribution of
warps across the MFMA tiles looked like:

w0 w1 w0 w1
w2 w3 w2 w3
w0 w1 w0 w1
w2 w3 w2 w3

The new tilesPerWarp parameter allows each warp to compute contiguous
MFMA tiles
in the row and/or column dimensions. Using the same example with
tilesPerWarp = [2, 2], the layout becomes:

w0 w0 w1 w1
w0 w0 w1 w1
w2 w2 w3 w3
w2 w2 w3 w3

While this is a general enhancement, the main motivation for introducing
this parameter
is to improve memory access efficiency for scale tensors in scaled dot
operations.
Specific patterns and use cases will be implemented in follow-up PRs.

---------

Co-authored-by: Ognjen Plavsic <[email protected]>
Co-authored-by: Lei Zhang <[email protected]>

* [AMD] Add support for pingpong GEMM using async_copy

* [BACKEND] combineRedundantWaitOps should not combine across loops/branches (triton-lang#7593)

`combineRedundantWaitOps` did skip over branches/loops, so if we end up
with something like:

```mlir
ttg.async_wait
scf.for
  ....
  scf.yield
ttg.async_wait
```
we merge the async_waits in the prologue and epilogue because we do not
find a `ttg.commit_group` in between. This PR stops the forward search
if we encounter a branch/loop. I can also walk through all successor
blocks if we think this is worth the effort.

This problem was not triggered before because the `ttg.async_wait` was
scheduled in the same stage as its user(s) so we ended up with no
`ttg.async_wait` in the prologue or there was another prefetch after it
in the prologue.

Since triton-lang#7458 we might place the
`ttg.async_wait` in the previous stage compared to its user(s) so we
might end up with the problematic IR.

* [AMD][NFC] Group scheduling functions in StreamPipeliner (triton-lang#7607)

NFC: Groups all scheduling related function to a namespace to prepare
for additional scheduling variants.

* [AMD] Add pingpong transformation for chained dot schedule (triton-lang#7638)

Adds support to enable pingpong for loops scheduled with the new
`ChainedDotSchedule` introduced by
triton-lang#7601.

The schedule already places the ops in the correct order so we just have
to insert the sync ops to ensure proper pingpong'ing.

* [AMD] Fix pingpong ChainedDot for empty second memory cluster (triton-lang#7694)

triton-lang#7638 introduced a null
pointer access (during review adjustments) if the second memory cluster
is empty or if there are no memory clusters at all.

Added a lit test to catch it and revert to the old logic.

* [AMD] Remove bypass permute optimization for AsyncCopy (triton-lang#7704)

We can only bypass ds_bpermute to apply the swizzling if lanes loading
the same row read a contiguous chunk of memory from HBM, which we cannot
infer when lowering to LLVM.
The current selection does only check if the elements for each lane are
contiguous which is not strict enough.

* [AMD] Add ChainedDotSchedule to StreamPipeliner (triton-lang#7601)

Adds a new scheduling variant which kicks in for loop which have 2
chained dots and `num_stages==4`. It places the two dots in consecutive
stages so we can interleave operations using the result of the first dot
with both dots in the loop, a pseudo example IR:

```
   %1 = tt.dot ...
   %2 = arith.addf %1, %arg1
   %3 = arith.subf %2, %arg2
   %4 = tt.dot %X, %Y, %3
```
Which could result in the following pseudo schedule (ignoring mem ops)
to interleave with both dots:
```
   stage N,   Cluster0: [%1 = tt.dot, %3 = arith.subf]
   stage N+1, Cluster1: [%4 = tt.dot, %2 = arith.addf]
```

As a first step the schedule splits the op chain between dot1 and dot2
when it encounters an operation which has more than 2 users. This aims
to avoid adding too many loop carried dependencies but does not
guarantee a good work balance between the two clusters. In future PRs we
might make this more sophisticated.

* [AMD] Add scale preshuffling and opSel implementation (triton-lang#7603)

This PR implements test kernel for efficient scale packing for CDNA4
arch as well as opSel for scaled MFMA instructions.

Scaled MFMA instructions expect scale operands as 32-bit values,
even though each individual scale is only 8 bits.
To reduce register usage, we pack 4 scales into a single 32-bit value
and use the opSel field to select the appropriate byte during execution.
Packing is done along the K dimension first. if there aren’t enough
values in K, we continue along the non-K dimension.

---------

Co-authored-by: Ognjen Plavsic <[email protected]>

* [AMD] Enable Pingpong by default on gfx950 arch (triton-lang#7697)

List of enabling conditions
- FP/BF16 GEMM with M,N>64 tilesize when num_stages=3 and num_warps=8
- GEMM using `dot_scaled` with M=N=256 tile size when num_stages=2 and
num_warps=8
- FA with num_stages=4
Only with using async_copy.

* [Backend] Bump to llvm/llvm-project@570885128351 (triton-lang#7291)

This picks up a bug fix for AMDGPU v_permlane_swap:

llvm/llvm-project#144423
Without this fix, the v_permlane_swap is wrongly sunk.

Along the way we need to fix API changes:

Add header file for the class IRBuilder
Add missing default parameter in convertFuncOpToLLVMFuncOp

---------

Co-authored-by: Maksim Levental <[email protected]>
Co-authored-by: Yi Qian <[email protected]>
Co-authored-by: Lei Zhang <[email protected]>
Co-authored-by: Lixun Zhang <[email protected]>
Co-authored-by: Jungwook Park <[email protected]>
Co-authored-by: Pengzhan Zhao <[email protected]>
Co-authored-by: plognjen <[email protected]>
Co-authored-by: Ognjen Plavsic <[email protected]>
Co-authored-by: Zeng Wu <[email protected]>

Author: 10 people

cmake/llvm-hash.txt

@@ -1 +1 @@
-8957e64a20fc7f4277565c6cfe3e555c119783ce
+570885128351868c1308bb22e8ca351d318bc4a1

include/triton/Conversion/TritonGPUToLLVM/TargetInfoBase.h

@@ -97,11 +97,6 @@ class TargetInfoBase {
   virtual bool supportLdMatrix() const { return false; }
   virtual bool supportStMatrix() const { return false; }
 
-  // Annotate target specific information to local store operations during
-  // lowering to LLVM.
-  virtual void localStoreOpAnnotation(triton::gpu::LocalStoreOp op,
-                                      size_t localStoreOpCount,
-                                      Type type) const {}
   // Annotate target specific information to local load operations during
   // lowering to LLVM. `llLoadOp` is the generated LLVM load op.
   virtual void localLoadOpAnnotation(triton::gpu::LocalLoadOp localLoadOp,

include/triton/Conversion/TritonGPUToLLVM/Utility.h

@@ -556,11 +556,12 @@ SmallVector<Value> loadSharedToDistributed(triton::gpu::LocalLoadOp localLoadOp,
                                            Location loc, RewriterBase &rewriter,
                                            const TargetInfoBase &target);
 
-void storeDistributedToShared(
-    triton::gpu::MemDescType dstTy, RankedTensorType srcTy, Type elemLlvmTy,
-    ArrayRef<Value> srcVals, const SharedMemoryObject &smemObj, Location loc,
-    RewriterBase &rewriter, const TargetInfoBase &target,
-    std::pair<size_t, Type> *const llvmOpCount = nullptr);
+void storeDistributedToShared(triton::gpu::MemDescType dstTy,
+                              RankedTensorType srcTy, Type elemLlvmTy,
+                              ArrayRef<Value> srcVals,
+                              const SharedMemoryObject &smemObj, Location loc,
+                              RewriterBase &rewriter,
+                              const TargetInfoBase &target);
 
 SmallVector<Value> unpackLLElements(Location loc, Value llvmStruct,
                                     RewriterBase &rewriter);

include/triton/Dialect/TritonGPU/IR/LinearLayoutConversions.h

@@ -282,10 +282,11 @@ LinearLayout getTmemLoadLayoutSplitLongM(int M, int N, RankedTensorType oldType,
                                          int numWarps);
 
 // Create LinearLayout for scale in scaled mfma.
-LinearLayout chooseScaledMfmaScaleLayout(
-    MLIRContext *ctx, int dotOperandIdx,
-    const std::vector<std::vector<int32_t>> &dotOperandWarpBasis,
-    ArrayRef<int64_t> dotOperandShape, unsigned mfmaMDim);
+LinearLayout chooseScaledMfmaScaleLayout(MLIRContext *ctx, int dotOperandIdx,
+                                         ArrayRef<int64_t> dotOperandShape,
+                                         unsigned mfmaMDim,
+                                         ArrayRef<unsigned> tilesPerWarp,
+                                         ArrayRef<unsigned> warpsPerCTA);
 
 // Create LinearLayout for nvidia mma tile.
 LinearLayout nvidiaMmaTile(MLIRContext *ctx, ArrayRef<unsigned> tileShape,

include/triton/Dialect/TritonGPU/IR/TritonGPUAttrDefs.td

@@ -919,11 +919,11 @@ An encoding for tensors that have been produced by MFMA matrix core instructions
 available on AMD Instinct GPUs of CDNA architectures.
 
 It is characterized by the following parameters:
-- `versionMajor` and `versionMinor` indicates the GPU architecture:
-  - 1.0: gfx908, i.e. CDNA1
-  - 2.0: gfx90a: i.e. CDNA2
-  - 3.0: gfx942: CDNA3
-  - 4.0: gfx950: CDNA4
+- `version` indicates the GPU architecture:
+  - 1: gfx908: CDNA1
+  - 2: gfx90a: CDNA2
+  - 3: gfx942: CDNA3
+  - 4: gfx950: CDNA4
 - `warpsPerCTA` indicates the warp layout in the block.
 - `MDim` and `NDim` indicate the dimension of the output of the mfma instruction.
 - `isTransposed` indicates the result tensor is transposed so that it can be converted to dotOperand layout
@@ -1009,24 +1009,61 @@ V [ 0,4,8...60   1,5...61     2,6...62     3,7...63    ]   [ 128,132...188  129,
   [ 64,68...124  65,69...125  66,70...126  67,71...127 ]   [ 192,196...252  193,197...253  194,198...254  195,199...255 ]
   [ 64,68...124  65,69...125  66,70...126  67,71...127 ]   [ 192,196...252  193,197...253  194,198...254  195,199...255 ]
   [ 64,68...124  65,69...125  66,70...126  67,71...127 ]   [ 192,196...252  193,197...253  194,198...254  195,199...255 ]
+
+Example 4:
+This example demonstrates semantics of tilesPerWarp parameter. The MFMA layout (with tilesPerWarp=[1,1])
+assumes that each warp within a CTA tile computes a single MFMA tile. When the tensor is larger than
+a single CTA tile, these tiles are repeated across the tensor. In this setup, the output tiles computed
+by each wave were strided by the number of warps per CTA tile in both row and column dimensions.
+
+For instance, with 16 MFMA tiles and warpsPerCTA = [2, 2], the distribution of warps across the MFMA
+tiles looked like:
+
+w0 w1 w0 w1
+w2 w3 w2 w3
+w0 w1 w0 w1
+w2 w3 w2 w3
+
+tilesPerWarp parameter allows each warp to compute contiguous MFMA tiles in the row and/or column dimensions.
+Using the same example with tilesPerWarp = [2, 2], the layout becomes:
+
+w0 w0 w1 w1
+w0 w0 w1 w1
+w2 w2 w3 w3
+w2 w2 w3 w3
 }];
 
   let parameters = (
     ins
-    "unsigned": $versionMajor,
-    "unsigned": $versionMinor,
+    "unsigned": $version,
     ArrayRefParameter<"unsigned">:$warpsPerCTA,
+    ArrayRefParameter<"unsigned">:$tilesPerWarp,
     "unsigned":$MDim,
     "unsigned":$NDim,
     "bool":$isTransposed,
     "CTALayoutAttr":$CTALayout
   );
 
+  let builders = [
+    AttrBuilder<(ins "unsigned":$version,
+                     "ArrayRef<unsigned>":$warpsPerCTA,
+                     "unsigned":$MDim,
+                     "unsigned":$NDim,
+                     "bool":$isTransposed,
+                     "CTALayoutAttr":$CTALayout), [{
+      SmallVector<unsigned> tilesPerWarp(warpsPerCTA.size(), 1);
+      return $_get(context, version, warpsPerCTA, tilesPerWarp, MDim, NDim, isTransposed, CTALayout);
+    }]>
+  ];
+
   let extraClassDeclaration = extraDistributedDeclaration # [{
     SmallVector<int64_t> getInstrShapeForOperand(int kWidth, int opIdx) const;
     SmallVector<int64_t> getRepForOperand(ArrayRef<int64_t> operandShape, int kWidth, int opIdx) const;
     SmallVector<unsigned> getRepOrderForOperand(int opIdx) const;
 
+    // Check if tilesPerWarp is 1 in every dimension.
+    bool hasUnitTilesPerWarp() const;
+
     // Returns a swizzled shared layout matching this MFMA layout for the
     // dot operand at the given |operandIdx| with |operandShape|.
     SwizzledSharedEncodingAttr composeSharedLayoutForOperand(

include/triton/Dialect/TritonGPU/Transforms/Schedule.h

@@ -4,6 +4,7 @@
 #include "mlir/Dialect/SCF/IR/SCF.h"
 #include "mlir/IR/ImplicitLocOpBuilder.h"
 #include "mlir/Support/LLVM.h"
+#include "triton/Analysis/AxisInfo.h"
 #include "triton/Dialect/TritonGPU/Transforms/PipelineExpander.h"
 #include "llvm/ADT/ArrayRef.h"
 #include <list>
@@ -17,6 +18,13 @@ namespace gpu {
 /// Lower the loops to prepare them for pipeline expansion.
 void lowerLoops(ModuleOp moduleOp);
 
+bool hasGpuBarriers(scf::ForOp forOp);
+bool isSafeToPipeline(scf::ForOp forOp);
+llvm::MapVector<Operation *, std::pair<int, Operation *>>
+loadOpsToIndirectionLevel(scf::ForOp forOp, bool pipelineWithoutDot,
+                          triton::ModuleAxisInfoAnalysis &axisInfoAnalysis,
+                          int numStages, bool filterSmall = true);
+
 }; // namespace gpu
 
 /// Pipeline the TMA stores in the loop.
@@ -191,6 +199,13 @@ class OpBuilderForStage : public mlir::ImplicitLocOpBuilder,
   CoarseSchedule &schedule;
 };
 
+namespace gpu {
+void scheduleDistanceOneDependencies(scf::ForOp forOp,
+                                     CoarseSchedule &schedule);
+void scheduleRemainingToLastStage(scf::ForOp forOp, CoarseSchedule &schedule,
+                                  CoarseSchedule::Cluster afterPrologue);
+} // namespace gpu
+
 } // namespace triton
 } // namespace mlir
 #endif // TRITON_TRITONGPU_TRANSFORM_PIPELINE_SCHEDULE_H_

include/triton/Dialect/TritonGPU/Transforms/Utility.h

@@ -206,7 +206,7 @@ bool isPureUnaryInlineAsm(Operation *op);
 int getNVIDIAComputeCapability(Operation *module);
 
 // Read the amd target from the module attributes
-StringRef getAMDArch(Operation *module);
+std::optional<StringRef> getAMDArch(Operation *module);
 
 std::optional<mlir::triton::gpu::SwizzledSharedEncodingAttr>
 getSharedEncIfAllUsersAreDotEnc(Value val, bool &incompatible);
@@ -258,11 +258,12 @@ void replaceUsesAndPropagateType(OpBuilder &builder, Operation *oldUse,
 
 /// Replace all uses of `old` with a local load from `alloc` unless the use is a
 /// `ttg.local_alloc` with a matching shared encoding, in which case the shared
-/// memory is forwarded directly into the use.
-void replaceUsesWithLocalLoad(
-    OpBuilder &builder, OpResult old,
-    TypedValue<triton::gpu::MemDescType> alloc,
-    TypedValue<triton::gpu::AsyncTokenType> token = {});
+/// memory is forwarded directly into the use. Returns the `ttg.local_load` if
+/// it created one.
+triton::gpu::LocalLoadOp
+replaceUsesWithLocalLoad(OpBuilder &builder, OpResult old,
+                         TypedValue<triton::gpu::MemDescType> alloc,
+                         TypedValue<triton::gpu::AsyncTokenType> token = {});
 
 // Return true if the value comes from a load or a block argument.
 // This will skip convert layouts and memdesc views.

lib/Conversion/TritonGPUToLLVM/ConvertLayoutOpToLLVM.cpp

@@ -18,6 +18,7 @@ namespace {
 using namespace mlir;
 using namespace mlir::triton::gpu;
 
+constexpr int kPtrBitWidth = 64;
 struct ConvertLayoutOpUsingLinearLayoutsConversion
     : public ConvertOpToLLVMPattern<ConvertLayoutOp> {
   const TargetInfoBase &targetInfo;

lib/Conversion/TritonGPUToLLVM/FuncOpToLLVM.cpp

@@ -1,15 +1,9 @@
+#include "mlir/Conversion/FuncToLLVM/ConvertFuncToLLVM.h"
 #include "mlir/Dialect/LLVMIR/NVVMDialect.h"
 #include "mlir/IR/BuiltinAttributes.h"
 #include "triton/Conversion/TritonGPUToLLVM/PatternTritonGPUOpToLLVM.h"
 #include "triton/Conversion/TritonGPUToLLVM/Utility.h"
 
-namespace mlir {
-FailureOr<LLVM::LLVMFuncOp>
-convertFuncOpToLLVMFuncOp(FunctionOpInterface funcOp,
-                          ConversionPatternRewriter &rewriter,
-                          const LLVMTypeConverter &converter);
-}
-
 namespace {
 
 using namespace mlir;

lib/Conversion/TritonGPUToLLVM/MemoryOpToLLVM.cpp

@@ -15,18 +15,19 @@ using namespace mlir::triton::gpu;
 // blocked -> shared.
 // Swizzling in shared memory to avoid bank conflict. Normally used for
 // A/B operands of dots.
-void lowerDistributedToShared(
-    Location loc, Value src, Value dst, Value adaptorSrc,
-    const SharedMemoryObject &smemObj, const LLVMTypeConverter *typeConverter,
-    ConversionPatternRewriter &rewriter, const TargetInfoBase &targetInfo,
-    std::pair<size_t, Type> *const llvmOpCount = nullptr) {
+void lowerDistributedToShared(Location loc, Value src, Value dst,
+                              Value adaptorSrc,
+                              const SharedMemoryObject &smemObj,
+                              const LLVMTypeConverter *typeConverter,
+                              ConversionPatternRewriter &rewriter,
+                              const TargetInfoBase &targetInfo) {
   auto srcTy = cast<RankedTensorType>(src.getType());
   auto dstTy = cast<MemDescType>(dst.getType());
   auto elemTy = typeConverter->convertType(srcTy.getElementType());
 
   auto inVals = unpackLLElements(loc, adaptorSrc, rewriter);
   storeDistributedToShared(dstTy, srcTy, elemTy, inVals, smemObj, loc, rewriter,
-                           targetInfo, llvmOpCount);
+                           targetInfo);
 }
 
 struct GlobalScratchAllocOpConversion
@@ -173,13 +174,9 @@ struct LocalStoreOpConversion
     auto smemObj = LLVM::getSharedMemoryObjectFromStruct(
         op.getLoc(), adaptor.getDst(), llvmElemTy, rewriter);
 
-    std::pair<size_t, Type> llvmOpCount;
     lowerDistributedToShared(op.getLoc(), op.getSrc(), op.getDst(),
                              adaptor.getSrc(), smemObj, getTypeConverter(),
-                             rewriter, targetInfo, &llvmOpCount);
-
-    targetInfo.localStoreOpAnnotation(op, llvmOpCount.first,
-                                      llvmOpCount.second);
+                             rewriter, targetInfo);
 
     rewriter.eraseOp(op);
     return success();