Integrate LLVM at 41f65666f6378bba7266be7c662c70074f04ed75

.gitmodules

@@ -1,6 +1,6 @@
 [submodule "externals/llvm-project"]
 	path = externals/llvm-project
-	url = https://github.com/llvm/llvm-project.git
+	url = https://github.com/iree-org/llvm-project.git
 [submodule "externals/stablehlo"]
 	path = externals/stablehlo
 	url = https://github.com/openxla/stablehlo.git

include/torch-mlir/Conversion/TorchToTosa/TosaLegalizeUtils.h

@@ -11,6 +11,7 @@
 #define TORCHMLIR_CONVERSION_TORCHTOTOSA_TOSALEGALIZEUTILS_H
 
 #include "mlir/Dialect/Quant/IR/QuantTypes.h"        // from @llvm-project
+#include "mlir/Dialect/Tosa/IR/TosaOps.h"            // from @llvm-project
 #include "mlir/Dialect/Tosa/Utils/ConversionUtils.h" // from @llvm-project
 #include "mlir/Dialect/Tosa/Utils/ShapeUtils.h"      // from @llvm-project
 #include "mlir/IR/BuiltinAttributes.h"               // from @llvm-project
@@ -26,8 +27,8 @@ namespace tosa {
 // rounding mode
 Value buildRescale(PatternRewriter &rewriter, Operation *op,
                    ShapedType output_type, Value input_val, double scale,
-                   int64_t input_zp, int64_t output_zp, StringRef rounding_mode,
-                   bool scale32);
+                   int64_t input_zp, int64_t output_zp,
+                   tosa::RoundingMode rounding_mode, bool scale32);
 
 // Creates TOSA rescale op with int32 output
 Value buildRescaleToInt32(PatternRewriter &rewriter, Operation *op,

lib/Conversion/TorchToTosa/TorchToTosa.cpp

@@ -138,7 +138,9 @@ class ConvertAtenBinaryOp : public OpConversionPattern<AtenOpT> {
       // tosa.minimum
       binaryOp = rewriter.create<TosaOpT>(
           op->getLoc(), outTy, lhs, rhs,
-          /*nan_mode=*/rewriter.getStringAttr("PROPAGATE"));
+          /*nan_mode=*/
+          tosa::NanPropagationModeAttr::get(
+              rewriter.getContext(), tosa::NanPropagationMode::PROPAGATE));
     } else {
       binaryOp =
           tosa::createBinaryOpAndCast<TosaOpT>(rewriter, op, outTy, lhs, rhs);
@@ -907,7 +909,9 @@ LogicalResult ConvertAtenOp<AtenReluOp>::matchAndRewrite(
   // Use default NaN Propagation mode "PROPAGATE" for tosa.clamp
   rewriter.replaceOpWithNewOp<tosa::ClampOp>(
       op, outTy, self, minFloatAttr, maxFloatAttr,
-      /*nan_mode=*/rewriter.getStringAttr("PROPAGATE"));
+      /*nan_mode=*/
+      tosa::NanPropagationModeAttr::get(rewriter.getContext(),
+                                        tosa::NanPropagationMode::PROPAGATE));
   return success();
 }
 
@@ -1237,7 +1241,9 @@ LogicalResult ConvertAtenOp<AtenArgmaxOp>::matchAndRewrite(
         .create<tosa::ArgMaxOp>(
             op->getLoc(), getTypeConverter()->convertType(outputReduceTy),
             input, reduceDimAttr,
-            /*nan_mode=*/rewriter.getStringAttr("PROPAGATE"))
+            /*nan_mode=*/
+            tosa::NanPropagationModeAttr::get(
+                rewriter.getContext(), tosa::NanPropagationMode::PROPAGATE))
         .getResult();
   };
 
@@ -3925,7 +3931,9 @@ class ConvertAtenMinMaxDimOp : public OpConversionPattern<AtenOpT> {
           op->getLoc(),
           RankedTensorType::get(makeShapeLLVMCompatible(reducedShape),
                                 selfElemType),
-          self, dimAttr, /*nan_mode=*/rewriter.getStringAttr("PROPAGATE"));
+          self, dimAttr, /*nan_mode=*/
+          tosa::NanPropagationModeAttr::get(
+              rewriter.getContext(), tosa::NanPropagationMode::PROPAGATE));
     } else {
       reduceOp = rewriter.create<TosaOpT>(
           op->getLoc(),
@@ -3946,14 +3954,18 @@ class ConvertAtenMinMaxDimOp : public OpConversionPattern<AtenOpT> {
           op->getLoc(),
           RankedTensorType::get(makeShapeLLVMCompatible(prunedShape),
                                 indicesElemType),
-          negateOp, dimAttr, /*nan_mode=*/rewriter.getStringAttr("PROPAGATE"));
+          negateOp, dimAttr, /*nan_mode=*/
+          tosa::NanPropagationModeAttr::get(
+              rewriter.getContext(), tosa::NanPropagationMode::PROPAGATE));
     } else {
       // Use default NaN Propagation mode "PROPAGATE" for tosa.argmax
       argMaxOp = rewriter.create<tosa::ArgMaxOp>(
           op->getLoc(),
           RankedTensorType::get(makeShapeLLVMCompatible(prunedShape),
                                 indicesElemType),
-          self, dimAttr, /*nan_mode=*/rewriter.getStringAttr("PROPAGATE"));
+          self, dimAttr, /*nan_mode=*/
+          tosa::NanPropagationModeAttr::get(
+              rewriter.getContext(), tosa::NanPropagationMode::PROPAGATE));
     }
 
     if (argMaxOp.getType() != indicesType) {
@@ -5202,7 +5214,9 @@ LogicalResult ConvertAtenOp<AtenClampOp>::matchAndRewrite(
 
     rewriter.replaceOpWithNewOp<tosa::ClampOp>(
         op, outType, adaptor.getSelf(), minIntAttr, maxIntAttr,
-        /*nan_mode=*/rewriter.getStringAttr("PROPAGATE"));
+        /*nan_mode=*/
+        tosa::NanPropagationModeAttr::get(rewriter.getContext(),
+                                          tosa::NanPropagationMode::PROPAGATE));
   } else {
     FloatAttr minFloatAttr, maxFloatAttr;
     if (outElemTy.isF16()) {
@@ -5231,7 +5245,9 @@ LogicalResult ConvertAtenOp<AtenClampOp>::matchAndRewrite(
 
     rewriter.replaceOpWithNewOp<tosa::ClampOp>(
         op, outType, adaptor.getSelf(), minFloatAttr, maxFloatAttr,
-        /*nan_mode=*/rewriter.getStringAttr("PROPAGATE"));
+        /*nan_mode=*/
+        tosa::NanPropagationModeAttr::get(rewriter.getContext(),
+                                          tosa::NanPropagationMode::PROPAGATE));
   }
 
   return success();
@@ -5340,13 +5356,17 @@ LogicalResult ConvertAtenOp<AtenClampTensorOp>::matchAndRewrite(
   // Use default NaN Propagation mode "PROPAGATE" for tosa.maximum
   auto minThresholdCheck = rewriter.create<tosa::MaximumOp>(
       op->getLoc(), resultType, self, min,
-      /*nan_mode=*/rewriter.getStringAttr("PROPAGATE"));
+      /*nan_mode=*/
+      tosa::NanPropagationModeAttr::get(rewriter.getContext(),
+                                        tosa::NanPropagationMode::PROPAGATE));
 
   // yi = min(max(xi, min_valuei), max_valuei)
   // Use default NaN Propagation mode "PROPAGATE" for tosa.minimum
   auto result = rewriter.create<tosa::MinimumOp>(
       op->getLoc(), resultType, minThresholdCheck, max,
-      /*nan_mode=*/rewriter.getStringAttr("PROPAGATE"));
+      /*nan_mode=*/
+      tosa::NanPropagationModeAttr::get(rewriter.getContext(),
+                                        tosa::NanPropagationMode::PROPAGATE));
 
   rewriter.replaceOp(op, result);
   return success();
@@ -5934,7 +5954,10 @@ class ConvertAtenPoolingBaseOp : public OpConversionPattern<AtenOpT> {
       pooledOutput = rewriter
                          .create<TosaOpT>(
                              op->getLoc(), outputTy, input, kernel, stride, pad,
-                             /*nan_mode=*/rewriter.getStringAttr("PROPAGATE"))
+                             /*nan_mode=*/
+                             tosa::NanPropagationModeAttr::get(
+                                 rewriter.getContext(),
+                                 tosa::NanPropagationMode::PROPAGATE))
                          .getResult();
     } else if constexpr (std::is_same<TosaOpT, tosa::AvgPool2dOp>::value) {
       TypeAttr accType;
@@ -6830,11 +6853,11 @@ ConvertAtenOp<Aten__InterpolateSizeListScaleListOp>::matchAndRewrite(
     return rewriter.notifyMatchFailure(
         op, "Only nearest and bilinear interpolation modes supported");
 
-  std::string mode;
+  tosa::ResizeMode mode;
   if (pyMode == "bilinear") {
-    mode = "BILINEAR";
+    mode = tosa::ResizeMode::BILINEAR;
   } else {
-    mode = "NEAREST_NEIGHBOR";
+    mode = tosa::ResizeMode::NEAREST_NEIGHBOR;
   }
 
   bool alignCorners;
@@ -6896,7 +6919,7 @@ ConvertAtenOp<Aten__InterpolateSizeListScaleListOp>::matchAndRewrite(
     offset = 0;
 
     // If nearest neighbours we need to guarantee we round up.
-    if (mode == "NEAREST_NEIGHBOR" && alignCorners) {
+    if (mode == tosa::ResizeMode::NEAREST_NEIGHBOR && alignCorners) {
       offset += n / 2;
     }
 
@@ -6916,7 +6939,8 @@ ConvertAtenOp<Aten__InterpolateSizeListScaleListOp>::matchAndRewrite(
       tosa::getTosaConstShape(rewriter, op->getLoc(), {offset_y, offset_x});
   auto border =
       tosa::getTosaConstShape(rewriter, op->getLoc(), {border_y, border_x});
-  StringAttr modeAttr = rewriter.getStringAttr(mode);
+
+  auto modeAttr = tosa::ResizeModeAttr::get(rewriter.getContext(), mode);
 
   auto resizeOpResult =
       rewriter
@@ -8610,11 +8634,14 @@ LogicalResult ConvertAtenOp<AtenLogitOp>::matchAndRewrite(
   // Clamp input to [eps, 1 - eps] when eps is not None
   // Use default NaN Propagation mode "PROPAGATE" for tosa.clamp
   if (!isEpsNone) {
-    zi = rewriter
-             .create<tosa::ClampOp>(
-                 op->getLoc(), resultType, self, minFloatAttr, maxFloatAttr,
-                 /*nan_mode=*/rewriter.getStringAttr("PROPAGATE"))
-             .getResult();
+    zi =
+        rewriter
+            .create<tosa::ClampOp>(
+                op->getLoc(), resultType, self, minFloatAttr, maxFloatAttr,
+                /*nan_mode=*/
+                tosa::NanPropagationModeAttr::get(
+                    rewriter.getContext(), tosa::NanPropagationMode::PROPAGATE))
+            .getResult();
   }
 
   auto one =

lib/Conversion/TorchToTosa/TosaLegalizeCommon.cpp

@@ -8,6 +8,7 @@
 //===----------------------------------------------------------------------===//
 
 #include "torch-mlir/Conversion/TorchToTosa/TosaLegalizeCommon.h"
+#include "mlir/Dialect/Tosa/IR/TosaOps.h" // from @llvm-project
 #include "mlir/Dialect/Tosa/Utils/ConversionUtils.h"
 #include "torch-mlir/Conversion/Utils/Utils.h"
 #include "torch-mlir/Dialect/Torch/IR/TorchOps.h"
@@ -764,7 +765,9 @@ std::optional<Value> convertReduceOpCommon(
         // and tosa.reduce_max
         reduce_op = CreateOpAndInfer<T>(
             rewriter, op->getLoc(), reduce_type, val, axis_attr,
-            /*nan_mode=*/rewriter.getStringAttr("PROPAGATE"));
+            /*nan_mode=*/
+            tosa::NanPropagationModeAttr::get(
+                rewriter.getContext(), tosa::NanPropagationMode::PROPAGATE));
       } else {
         reduce_op = CreateOpAndInfer<T>(rewriter, op->getLoc(), reduce_type,
                                         val, axis_attr);
@@ -777,7 +780,7 @@ std::optional<Value> convertReduceOpCommon(
       RankedTensorType output_rescale_type =
           RankedTensorType::get(shape_vec, output_type.getElementType());
       val = buildRescale(rewriter, op, output_rescale_type, val, output_scale,
-                         0, output_zp, "SINGLE_ROUND", true);
+                         0, output_zp, tosa::RoundingMode::SINGLE_ROUND, true);
     }
 
     // Optionally squeeze out the reduced axes.

lib/Conversion/TorchToTosa/TosaLegalizeUtils.cpp

@@ -35,8 +35,8 @@ Value buildRescaleMultiplier(bool scale32, PatternRewriter &rewriter,
 // rounding mode
 Value buildRescale(PatternRewriter &rewriter, Operation *op,
                    ShapedType output_type, Value input_val, double scale,
-                   int64_t input_zp, int64_t output_zp, StringRef rounding_mode,
-                   bool scale32) {
+                   int64_t input_zp, int64_t output_zp,
+                   tosa::RoundingMode rounding_mode, bool scale32) {
   int32_t multiplier;
   int32_t shift;
 
@@ -70,7 +70,8 @@ Value buildRescale(PatternRewriter &rewriter, Operation *op,
   auto rescale_op = CreateOpAndInfer<tosa::RescaleOp>(
       rewriter, op->getLoc(), output_type, input_val, multiplier_val, shift_val,
       input_zp_val.value(), output_zp_val.value(),
-      rewriter.getBoolAttr(scale32), rewriter.getStringAttr(rounding_mode),
+      rewriter.getBoolAttr(scale32),
+      tosa::RoundingModeAttr::get(rewriter.getContext(), rounding_mode),
       rewriter.getBoolAttr(false), rewriter.getBoolAttr(input_unsigned),
       rewriter.getBoolAttr(output_unsigned));
 
@@ -87,7 +88,7 @@ Value buildRescaleToInt32(PatternRewriter &rewriter, Operation *op,
   auto output_type = input_type.clone(rewriter.getI32Type());
 
   return buildRescale(rewriter, op, output_type, input_val, input_scale,
-                      input_zp, 0, "SINGLE_ROUND", true);
+                      input_zp, 0, tosa::RoundingMode::SINGLE_ROUND, true);
 }
 
 // Creates a TOSA rescale op based on conv2d parameters.
@@ -146,7 +147,9 @@ Value buildRescaleOpConvOutput(PatternRewriter &rewriter, Operation *op,
     auto rescale_op = CreateOpAndInfer<tosa::RescaleOp>(
         rewriter, op->getLoc(), output_type, conv_val, multiplier_val,
         shift_val, input_zp_val.value(), output_zp_val.value(),
-        rewriter.getBoolAttr(scale32), rewriter.getStringAttr("DOUBLE_ROUND"),
+        rewriter.getBoolAttr(scale32),
+        tosa::RoundingModeAttr::get(rewriter.getContext(),
+                                    tosa::RoundingMode::DOUBLE_ROUND),
         rewriter.getBoolAttr(false), rewriter.getBoolAttr(input_unsigned),
         rewriter.getBoolAttr(output_unsigned));
 
@@ -188,7 +191,9 @@ Value buildRescaleOpConvOutput(PatternRewriter &rewriter, Operation *op,
     auto rescale_op = CreateOpAndInfer<tosa::RescaleOp>(
         rewriter, op->getLoc(), output_type, conv_val, multiplier_val,
         shift_val, input_zp_val.value(), output_zp_val.value(),
-        rewriter.getBoolAttr(scale32), rewriter.getStringAttr("DOUBLE_ROUND"),
+        rewriter.getBoolAttr(scale32),
+        tosa::RoundingModeAttr::get(rewriter.getContext(),
+                                    tosa::RoundingMode::DOUBLE_ROUND),
         rewriter.getBoolAttr(true), rewriter.getBoolAttr(input_unsigned),
         rewriter.getBoolAttr(output_unsigned));
 

lib/Dialect/Torch/IR/TorchOps.cpp

@@ -373,21 +373,21 @@ LogicalResult ClassTypeOp::verify() {
 // PrimLoopOp
 //===----------------------------------------------------------------------===//
 
-OperandRange PrimLoopOp::getEntrySuccessorOperands(RegionBranchPoint point) {
-  assert(point == getRegion());
+OperandRange PrimLoopOp::getEntrySuccessorOperands(RegionSuccessor successor) {
+  assert(successor.getSuccessor() == &getRegion());
   return getIterArgsInit();
 }
 
 void PrimLoopOp::getSuccessorRegions(
     RegionBranchPoint point, SmallVectorImpl<RegionSuccessor> &regions) {
   Region &region = getRegion();
-  if (!point.getRegionOrNull()) {
+  if (!point.getTerminatorPredecessorOrNull()) {
     regions.emplace_back(&region, region.getArguments().slice(1));
     return;
   }
-  assert(point == region);
+  assert(point.getTerminatorPredecessorOrNull()->getParentRegion() == &region);
   regions.emplace_back(&region, region.getArguments().slice(1));
-  regions.emplace_back(getResults());
+  regions.emplace_back(getOperation(), getResults());
 }
 
 bool PrimLoopOp::isForLike() {
@@ -400,7 +400,7 @@ bool PrimLoopOp::isForLike() {
 //===----------------------------------------------------------------------===//
 
 MutableOperandRange
-PrimLoopConditionOp::getMutableSuccessorOperands(RegionBranchPoint point) {
+PrimLoopConditionOp::getMutableSuccessorOperands(RegionSuccessor successor) {
   // Pass all operands except the condition to the successor which is the
   // parent loop op.
   return getIterArgsMutable();
@@ -452,8 +452,8 @@ void PrimIfOp::print(OpAsmPrinter &p) {
 void PrimIfOp::getSuccessorRegions(RegionBranchPoint point,
                                    SmallVectorImpl<RegionSuccessor> &regions) {
   // The `then` and the `else` region branch back to the parent operation.
-  if (point.getRegionOrNull()) {
-    regions.push_back(RegionSuccessor(getResults()));
+  if (point.getTerminatorPredecessorOrNull()) {
+    regions.push_back(RegionSuccessor(getOperation(), getResults()));
     return;
   }
 
@@ -5321,17 +5321,18 @@ template <typename CalculateOp>
 static void
 getSuccessorRegionsForCalculateOp(CalculateOp op, RegionBranchPoint point,
                                   SmallVectorImpl<RegionSuccessor> &regions) {
-  if (!point.getRegionOrNull()) {
+  if (!point.getTerminatorPredecessorOrNull()) {
     // First thing the op does is branch into the calculation.
     regions.emplace_back(&op.getCalculation());
     return;
   }
-  if (point == op.getBody()) {
+  Region *region = point.getTerminatorPredecessorOrNull()->getParentRegion();
+  if (region == &op.getBody()) {
     // Body returns control to the outer op, passing through results.
-    regions.emplace_back(op.getResults());
+    regions.emplace_back(op.getOperation(), op.getResults());
     return;
   }
-  assert(point == op.getCalculation());
+  assert(region == &op.getCalculation());
   // Calculation branches to the body.
   regions.emplace_back(&op.getBody());
 }
@@ -5355,7 +5356,7 @@ void DtypeCalculateOp::getSuccessorRegions(
 //===----------------------------------------------------------------------===//
 
 MutableOperandRange ShapeCalculateYieldShapesOp::getMutableSuccessorOperands(
-    RegionBranchPoint point) {
+    RegionSuccessor successor) {
   // The shape operands don't get forwarded to the body.
   // MutableOperandRange always has an owning operation, even if empty, so
   // create a 0-length range.
@@ -5846,7 +5847,7 @@ LogicalResult AtenKthvalueOp::verify() {
 //===----------------------------------------------------------------------===//
 
 MutableOperandRange DtypeCalculateYieldDtypesOp::getMutableSuccessorOperands(
-    RegionBranchPoint point) {
+    RegionSuccessor successor) {
   // The dtype operands don't get forwarded to the body.
   // MutableOperandRange always has an owning operation, even if empty, so
   // create a 0-length range.

lib/RefBackend/RefBackend.cpp

@@ -361,7 +361,8 @@ class ExpandOpsForLLVM : public ExpandOpsForLLVMBase<ExpandOpsForLLVM> {
     auto func = getOperation();
     auto *context = &getContext();
     RewritePatternSet patterns(context);
-    populateExpandTanhPattern(patterns);
+    math::populateExpansionPatterns(patterns,
+                                    {math::TanhOp::getOperationName()});
     patterns.add<math::ErfPolynomialApproximation>(patterns.getContext());
     ConversionTarget target(*context);
     target.addLegalDialect<func::FuncDialect>();