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,7 @@ class ExpandOpsForLLVM : public ExpandOpsForLLVMBase<ExpandOpsForLLVM> {
  auto func = getOperation();
  auto *context = &getContext();
  RewritePatternSet patterns(context);
- populateExpandTanhPattern(patterns);
+ math::populateExpansionPatterns(patterns, {"tanh"});
  patterns.add<math::ErfPolynomialApproximation>(patterns.getContext());
  ConversionTarget target(*context);
  target.addLegalDialect<func::FuncDialect>();

projects/pt1/e2e_testing/xfail_sets.py

@@ -3481,13 +3481,9 @@
 if torch_version_for_comparison() > version.parse("2.4.0.dev"):
  STABLEHLO_PASS_SET = STABLEHLO_PASS_SET - {
  "ElementwiseCreateComplexModule_basic",
- "ElementwiseTanIntModule_basic",
- "ElementwiseTanModule_basic",
  }
  FX_IMPORTER_STABLEHLO_XFAIL_SET = FX_IMPORTER_STABLEHLO_XFAIL_SET | {
  "ElementwiseCreateComplexModule_basic",
- "ElementwiseTanIntModule_basic",
- "ElementwiseTanModule_basic",
  }