[mlir][NVGPU] Fix incorrect API usage in RewritePatterns

[springerm]

Summary

Incorrect API usage was detected by D144552.

mlir/lib/Dialect/NVGPU/Transforms/MmaSyncTF32Transform.cpp

struct MmaSyncF32ToTF32Pattern : public OpRewritePattern<nvgpu::MmaSyncOp> {
   using OpRewritePattern<nvgpu::MmaSyncOp>::OpRewritePattern;
 
   MmaSyncF32ToTF32Pattern(MLIRContext *context,
                           nvgpu::MmaSyncF32Lowering precision)
       : OpRewritePattern<nvgpu::MmaSyncOp>(context, /*benifit*/ 1),
         precision(precision) {}
 
   LogicalResult matchAndRewrite(nvgpu::MmaSyncOp op,
-                                PatternRewriter &rewrite) const override {
+                                PatternRewriter &rewriter) const override {
     Location location = op->getLoc();
 
     if (op->hasAttr(op.getTf32EnabledAttrName()) ||
         !op.getMatrixA().getType().cast<VectorType>().getElementType().isF32())
       return failure();
 
     if (precision == MmaSyncF32Lowering::Unkown)
       return emitError(location, "MmaSync F32-to-TF32 cannot be lowered with "
                                  "unknown precision level");
 
     if (precision == MmaSyncF32Lowering::TF32x3)
       return emitError(location, "TF32x3 is not supported at the moment "
                                  "for nvgpu.mma.sync on f32 datatype");
 
-    if (precision == MmaSyncF32Lowering::TF32)
-      op.setTf32EnabledAttr(rewrite.getUnitAttr());
+    if (precision == MmaSyncF32Lowering::TF32) {
+      rewriter.updateRootInPlace(
+          op, [&]() { op.setTf32EnabledAttr(rewriter.getUnitAttr()); });
+    }
 
     return success();
   }
 
 private:
   /// Precision for F32 Tensor Cores (TF32 or TF32x3)
   nvgpu::MmaSyncF32Lowering precision;
 };
 
 } // namespace
 
 void mlir::nvgpu::populateMmaSyncF32ToTF32Patterns(
     RewritePatternSet &patterns, nvgpu::MmaSyncF32Lowering precision) {
 
   patterns.add<MmaSyncF32ToTF32Pattern>(patterns.getContext(), precision);
 }