diff --git a/build_tools/bazel_to_cmake/bazel_to_cmake_targets.py b/build_tools/bazel_to_cmake/bazel_to_cmake_targets.py
index c91d04ea323b..71f7bc207148 100644
--- a/build_tools/bazel_to_cmake/bazel_to_cmake_targets.py
+++ b/build_tools/bazel_to_cmake/bazel_to_cmake_targets.py
@@ -114,6 +114,7 @@ def __init__(self, repo_map: Dict[str, str]):
             "MhloPasses",
         ],
         "@mlir-hlo//stablehlo:chlo_ops": ["ChloOps",],
+        "@mlir-hlo//stablehlo:stablehlo_ops": ["StablehloOps",],
         "@mlir-hlo//:stablehlo_legalize_to_hlo_pass": ["StablehloToMhlo",],
         "@mlir-hlo//stablehlo:broadcast_utils": ["StablehloBroadcastUtils",],
 
diff --git a/compiler/src/iree/compiler/InputConversion/CMakeLists.txt b/compiler/src/iree/compiler/InputConversion/CMakeLists.txt
index 305c05ed7d97..474215f2d60d 100644
--- a/compiler/src/iree/compiler/InputConversion/CMakeLists.txt
+++ b/compiler/src/iree/compiler/InputConversion/CMakeLists.txt
@@ -8,6 +8,7 @@ add_subdirectory(Common)
 
 if(IREE_INPUT_MHLO)
   add_subdirectory(MHLO)
+  add_subdirectory(StableHLO)
 endif()
 if(IREE_INPUT_TORCH)
   add_subdirectory(TMTensor)
diff --git a/compiler/src/iree/compiler/InputConversion/StableHLO/BUILD.bazel b/compiler/src/iree/compiler/InputConversion/StableHLO/BUILD.bazel
new file mode 100644
index 000000000000..7c6a8bf6e3d8
--- /dev/null
+++ b/compiler/src/iree/compiler/InputConversion/StableHLO/BUILD.bazel
@@ -0,0 +1,97 @@
+# Copyright 2023 The IREE Authors
+#
+# Licensed under the Apache License v2.0 with LLVM Exceptions.
+# See https://llvm.org/LICENSE.txt for license information.
+# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+load("//build_tools/bazel:build_defs.oss.bzl", "iree_compiler_cc_library", "iree_gentbl_cc_library")
+
+package(
+    default_visibility = ["//visibility:public"],
+    features = ["layering_check"],
+    licenses = ["notice"],  # Apache 2.0
+)
+
+iree_gentbl_cc_library(
+    name = "PassesIncGen",
+    tbl_outs = [
+        (
+            ["--gen-pass-decls"],
+            "Passes.h.inc",
+        ),
+    ],
+    tblgen = "@llvm-project//mlir:mlir-tblgen",
+    td_file = "Passes.td",
+    deps = [
+        "@llvm-project//mlir:PassBaseTdFiles",
+    ],
+)
+
+iree_compiler_cc_library(
+    name = "PassHeaders",
+    hdrs = [
+        "PassDetail.h",
+        "Passes.h",
+        "Passes.h.inc",
+        "Rewriters.h",
+        "TypeConversion.h",
+    ],
+    deps = [
+        ":PassesIncGen",
+        "@llvm-project//mlir:Pass",
+        "@llvm-project//mlir:Transforms",
+    ],
+)
+
+iree_compiler_cc_library(
+    name = "StableHLO",
+    srcs = [
+        "LegalizeToLinalgUtils.cpp",
+        "Passes.cpp",
+        "StableHLOToLinalg.cpp",
+        "TypeConversion.cpp",
+    ],
+    hdrs = [
+        "LegalizeToLinalgUtils.h",
+        "MapStableHLOToScalarOp.h",
+        "Passes.h",
+    ],
+    deps = [
+        ":PassHeaders",
+        ":PassesIncGen",
+        "@llvm-project//llvm:Support",
+        "@llvm-project//mlir:AffineDialect",
+        "@llvm-project//mlir:AffineUtils",
+        "@llvm-project//mlir:ArithDialect",
+        "@llvm-project//mlir:BufferizationDialect",
+        "@llvm-project//mlir:ComplexDialect",
+        "@llvm-project//mlir:ControlFlowDialect",
+        "@llvm-project//mlir:DialectUtils",
+        "@llvm-project//mlir:FuncDialect",
+        "@llvm-project//mlir:FuncTransforms",
+        "@llvm-project//mlir:IR",
+        "@llvm-project//mlir:LinalgDialect",
+        "@llvm-project//mlir:LinalgTransforms",
+        "@llvm-project//mlir:LinalgUtils",
+        "@llvm-project//mlir:MLProgramDialect",
+        "@llvm-project//mlir:MathDialect",
+        "@llvm-project//mlir:MemRefDialect",
+        "@llvm-project//mlir:Pass",
+        "@llvm-project//mlir:ReconcileUnrealizedCasts",
+        "@llvm-project//mlir:SCFDialect",
+        "@llvm-project//mlir:SCFToControlFlow",
+        "@llvm-project//mlir:SCFTransforms",
+        "@llvm-project//mlir:ShapeDialect",
+        "@llvm-project//mlir:ShapeToStandard",
+        "@llvm-project//mlir:ShapeTransforms",
+        "@llvm-project//mlir:SparseTensorDialect",
+        "@llvm-project//mlir:Support",
+        "@llvm-project//mlir:TensorDialect",
+        "@llvm-project//mlir:TensorUtils",
+        "@llvm-project//mlir:Transforms",
+        "@llvm-project//mlir:VectorDialect",
+        "@mlir-hlo//stablehlo:broadcast_utils",
+        "@mlir-hlo//stablehlo:chlo_ops",
+        "@mlir-hlo//stablehlo:stablehlo_ops",
+    ],
+)
diff --git a/compiler/src/iree/compiler/InputConversion/StableHLO/CMakeLists.txt b/compiler/src/iree/compiler/InputConversion/StableHLO/CMakeLists.txt
new file mode 100644
index 000000000000..78a3216e1487
--- /dev/null
+++ b/compiler/src/iree/compiler/InputConversion/StableHLO/CMakeLists.txt
@@ -0,0 +1,89 @@
+################################################################################
+# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from           #
+# compiler/src/iree/compiler/InputConversion/StableHLO/BUILD.bazel             #
+#                                                                              #
+# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary   #
+# CMake-only content.                                                          #
+#                                                                              #
+# To disable autogeneration for this file entirely, delete this header.        #
+################################################################################
+
+iree_add_all_subdirs()
+
+iree_tablegen_library(
+  NAME
+    PassesIncGen
+  TD_FILE
+    "Passes.td"
+  OUTS
+    --gen-pass-decls Passes.h.inc
+)
+
+iree_cc_library(
+  NAME
+    PassHeaders
+  HDRS
+    "PassDetail.h"
+    "Passes.h"
+    "Passes.h.inc"
+    "Rewriters.h"
+    "TypeConversion.h"
+  DEPS
+    ::PassesIncGen
+    MLIRPass
+    MLIRTransforms
+  PUBLIC
+)
+
+iree_cc_library(
+  NAME
+    StableHLO
+  HDRS
+    "LegalizeToLinalgUtils.h"
+    "MapStableHLOToScalarOp.h"
+    "Passes.h"
+  SRCS
+    "LegalizeToLinalgUtils.cpp"
+    "Passes.cpp"
+    "StableHLOToLinalg.cpp"
+    "TypeConversion.cpp"
+  DEPS
+    ::PassHeaders
+    ::PassesIncGen
+    ChloOps
+    LLVMSupport
+    MLIRAffineDialect
+    MLIRAffineUtils
+    MLIRArithDialect
+    MLIRBufferizationDialect
+    MLIRComplexDialect
+    MLIRControlFlowDialect
+    MLIRFuncDialect
+    MLIRFuncTransforms
+    MLIRIR
+    MLIRLinalgDialect
+    MLIRLinalgTransforms
+    MLIRLinalgUtils
+    MLIRMLProgramDialect
+    MLIRMathDialect
+    MLIRMemRefDialect
+    MLIRPass
+    MLIRReconcileUnrealizedCasts
+    MLIRSCFDialect
+    MLIRSCFToControlFlow
+    MLIRSCFTransforms
+    MLIRShapeDialect
+    MLIRShapeOpsTransforms
+    MLIRShapeToStandard
+    MLIRSparseTensorDialect
+    MLIRSupport
+    MLIRTensorDialect
+    MLIRTensorUtils
+    MLIRTransforms
+    MLIRVectorDialect
+    StablehloBroadcastUtils
+    StablehloOps
+  PUBLIC
+)
+
+### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ###
diff --git a/compiler/src/iree/compiler/InputConversion/StableHLO/LegalizeToLinalgUtils.cpp b/compiler/src/iree/compiler/InputConversion/StableHLO/LegalizeToLinalgUtils.cpp
new file mode 100644
index 000000000000..87f2f47de6c1
--- /dev/null
+++ b/compiler/src/iree/compiler/InputConversion/StableHLO/LegalizeToLinalgUtils.cpp
@@ -0,0 +1,132 @@
+// Copyright 2022 The IREE Authors
+//
+// Licensed under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+// Implements utilities for lowering StableHLO/CHLO dialect to Linalg dialect.
+
+#include "iree/compiler/InputConversion/StableHLO/LegalizeToLinalgUtils.h"
+
+#include <algorithm>
+#include <numeric>
+#include <string>
+#include <utility>
+
+#include "mlir/Dialect/Bufferization/IR/Bufferization.h"
+#include "mlir/Dialect/SparseTensor/IR/SparseTensor.h"
+#include "mlir/Dialect/Tensor/IR/Tensor.h"
+#include "stablehlo/dialect/ChloOps.h"
+#include "stablehlo/dialect/StablehloOps.h"
+
+namespace mlir::iree_compiler::stablehlo {
+namespace {
+bool hasIntegralShapeType(Operation* op) {
+  auto stp = op->getOperand(0).getType().dyn_cast<ShapedType>();
+  return stp && stp.getElementType().isIntOrIndex();
+}
+
+}  // namespace
+
+SmallVector<utils::IteratorType, 3> getParallelAndReductionIterators(
+    unsigned nLoops, unsigned nReduction) {
+  SmallVector<utils::IteratorType, 3> res(nLoops - nReduction,
+                                          utils::IteratorType::parallel);
+  res.append(nReduction, utils::IteratorType::reduction);
+  return res;
+}
+
+SmallVector<utils::IteratorType, 3> getNParallelLoopsAttrs(
+    unsigned nParallelLoops) {
+  return getParallelAndReductionIterators(nParallelLoops, 0);
+}
+
+Value getEmptySparseTensor(OpBuilder& b, Location loc, ShapedType type,
+                           ArrayRef<Value> dynSizes) {
+  return b.create<bufferization::AllocTensorOp>(loc, type.cast<TensorType>(),
+                                                dynSizes,
+                                                /*copy=*/Value(),
+                                                /*memory_space=*/IntegerAttr());
+}
+
+Value getEmptyTensor(OpBuilder& b, Location loc, ShapedType type,
+                     ArrayRef<Value> dynSizes) {
+  return b.create<tensor::EmptyOp>(loc, type.getShape(), type.getElementType(),
+                                   dynSizes,
+                                   type.cast<RankedTensorType>().getEncoding());
+}
+
+Value getEmptyTensorFor(OpBuilder& b, Location loc, ShapedType resultType,
+                        Operation* op, ValueRange operands) {
+  bool isSparse = sparse_tensor::getSparseTensorEncoding(resultType) != nullptr;
+  // Collect the sizes for a ranked tensor to be passed as parameter to a
+  // new tensor initialization operation. This operation only needs the
+  // dynamic sizes.
+  SmallVector<Value> sizes;
+  if (resultType.hasRank() && !resultType.hasStaticShape()) {
+    // Ask the op for its output shape.
+    auto shapeSource = cast<InferShapedTypeOpInterface>(op);
+    SmallVector<Value, 1> reifiedShapes;
+    (void)shapeSource.reifyReturnTypeShapes(b, operands, reifiedShapes);
+    assert(reifiedShapes.size() == 1 && "Expected one reified result");
+    // Construct sizes for the required dimensions.
+    for (const auto& en : llvm::enumerate(resultType.getShape())) {
+      if (en.value() != ShapedType::kDynamic) continue;
+      sizes.push_back(b.create<tensor::ExtractOp>(
+          loc, reifiedShapes[0],
+          ValueRange{b.create<arith::ConstantIndexOp>(loc, en.index())}));
+    }
+  }
+  return isSparse ? getEmptySparseTensor(b, loc, resultType, sizes)
+                  : getEmptyTensor(b, loc, resultType, sizes);
+}
+
+Value preSparsify(Operation* op, llvm::SmallVector<Value, 2>& values, Type rtp,
+                  OpBuilder* b) {
+  // Apply for semi-ring operations that lower to elaborate code
+  // (any sign-op, or an integral abs-op).
+  // TODO(peiming, ajcbik): these all can potentially be optimized by applying
+  // value transform on sparse_tenosr.value memref
+  if (isa<mlir::stablehlo::SignOp>(op) || isa<mlir::stablehlo::NegOp>(op) ||
+      (isa<mlir::stablehlo::AbsOp>(op) && hasIntegralShapeType(op)) ||
+      isa<chlo::AsinOp>(op) || isa<chlo::AsinhOp>(op) ||
+      isa<chlo::AtanOp>(op) || isa<chlo::AtanhOp>(op) ||
+      isa<chlo::BesselI1eOp>(op) || isa<chlo::SinhOp>(op) ||
+      isa<chlo::TanOp>(op)) {
+    if (!sparse_tensor::getSparseTensorEncoding(op->getResult(0).getType()) &&
+        !sparse_tensor::getSparseTensorEncoding(op->getOperand(0).getType()))
+      return Value();
+    Location loc = op->getLoc();
+    auto semiring = b->create<sparse_tensor::UnaryOp>(loc, rtp, values[0]);
+    Type itp = values[0].getType();
+    Block* present = b->createBlock(&semiring.getPresentRegion(), {}, itp, loc);
+    b->setInsertionPointToStart(&semiring.getPresentRegion().front());
+    values[0] = present->getArgument(0);
+    return semiring;
+  }
+  return Value();
+}
+
+Value postSparsify(Operation* op, Value semiring, Value result, OpBuilder* b) {
+  if (semiring) {
+    b->create<sparse_tensor::YieldOp>(op->getLoc(), result);
+    b->setInsertionPointAfter(semiring.getDefiningOp());
+    return semiring;
+  }
+  return result;
+}
+
+bool allOperandsAreScalarTensors(Operation* op) {
+  return llvm::all_of(op->getOperands(), [](Value operand) {
+    auto operandTy = operand.getType().dyn_cast<ShapedType>();
+    return operandTy && operandTy.getRank() == 0;
+  });
+}
+
+bool isInBodyOfLinalgOps(Operation* op) {
+  auto* parentOp = op->getParentRegion()->getParentOp();
+  return parentOp->getDialect() ==
+         parentOp->getContext()->getLoadedDialect<linalg::LinalgDialect>();
+}
+
+}  // namespace mlir::iree_compiler::stablehlo
diff --git a/compiler/src/iree/compiler/InputConversion/StableHLO/LegalizeToLinalgUtils.h b/compiler/src/iree/compiler/InputConversion/StableHLO/LegalizeToLinalgUtils.h
new file mode 100644
index 000000000000..5f36df166b9c
--- /dev/null
+++ b/compiler/src/iree/compiler/InputConversion/StableHLO/LegalizeToLinalgUtils.h
@@ -0,0 +1,178 @@
+// Copyright 2022 The IREE Authors
+//
+// Licensed under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+// Utils for lowering of the StableHLO dialect to the Linalg dialect.
+
+#ifndef IREE_COMPILER_INPUTCONVERSION_STABLEHLO_LEGALIZE_TO_LINALG_UTILS_H_
+#define IREE_COMPILER_INPUTCONVERSION_STABLEHLO_LEGALIZE_TO_LINALG_UTILS_H_
+
+#include <algorithm>
+#include <numeric>
+#include <optional>
+#include <string>
+#include <utility>
+
+#include "iree/compiler/InputConversion/StableHLO/MapStableHLOToScalarOp.h"
+#include "llvm/ADT/STLExtras.h"
+#include "llvm/ADT/SmallVector.h"
+#include "llvm/ADT/StringSet.h"
+#include "mlir/Dialect/Linalg/IR/Linalg.h"
+#include "mlir/Dialect/Linalg/Utils/Utils.h"
+#include "mlir/Dialect/Tensor/IR/Tensor.h"
+#include "mlir/IR/AffineExpr.h"
+#include "mlir/IR/Attributes.h"
+#include "mlir/IR/Builders.h"
+#include "mlir/IR/BuiltinAttributes.h"
+#include "mlir/IR/BuiltinTypes.h"
+#include "mlir/IR/Location.h"
+#include "mlir/IR/MLIRContext.h"
+#include "mlir/IR/Operation.h"
+#include "mlir/IR/OperationSupport.h"
+#include "mlir/IR/TypeUtilities.h"
+#include "mlir/Support/LLVM.h"
+#include "mlir/Support/LogicalResult.h"
+#include "mlir/Transforms/DialectConversion.h"
+#include "stablehlo/dialect/StablehloOps.h"
+
+namespace mlir::iree_compiler::stablehlo {
+
+/// Returns an ArrayAttr that contains `nLoops` attributes. All the attributes
+/// are "parallel" except the last `nReduction` elements, where are "reduction"
+/// attributes.
+SmallVector<utils::IteratorType, 3> getParallelAndReductionIterators(
+    unsigned nLoops, unsigned nReduction);
+
+/// Returns an ArrayAttr that contains `nParallelLoops` "parallel" attributes.
+SmallVector<utils::IteratorType, 3> getNParallelLoopsAttrs(
+    unsigned nParallelLoops);
+
+/// Generates an init sparse tensor.
+Value getEmptySparseTensor(OpBuilder& b, Location loc, ShapedType type,
+                           ArrayRef<Value> dynSizes);
+
+/// Generates a tensor.empty op.
+Value getEmptyTensor(OpBuilder& b, Location loc, ShapedType type,
+                     ArrayRef<Value> dynSizes);
+
+/// Generates an empty tensor for the result of the operation, which could be a
+/// dense tensor or a sparse tensor.
+Value getEmptyTensorFor(OpBuilder& b, Location loc, ShapedType resultType,
+                        Operation* op, ValueRange operands);
+
+/// Sparsifies a (block of) operation(s) that cannot be handled directly
+/// by the sparse compiler but has well-known semi-ring semantics.
+///
+/// This yields something of the following form:
+///
+///   %result = sparse_tensor.unary %values[0]
+///     present={
+///       ^bb1(%val):
+///         ... codegen proceeds here using %val ....
+///         sparse_tensor.yield
+///     }
+///     absent={}
+///   linalg.yield %result
+Value preSparsify(Operation* op, llvm::SmallVector<Value, 2>& values, Type rtp,
+                  OpBuilder* b);
+
+/// Finalizes sparse semi-ring construction.
+Value postSparsify(Operation* op, Value semiring, Value result, OpBuilder* b);
+
+/// Returns true if all operands are tensors with rank 0.
+bool allOperandsAreScalarTensors(Operation* op);
+
+/// Returns true if parent op is linalg.
+bool isInBodyOfLinalgOps(Operation* op);
+
+/// Converts a HLO operation to a linalg.generic op that contains the
+/// corresponding scalar operations.
+template <typename OpTy>
+class PointwiseToLinalgConverter : public OpConversionPattern<OpTy> {
+ public:
+  using OpConversionPattern<OpTy>::OpConversionPattern;
+
+  LogicalResult matchAndRewrite(
+      OpTy op, typename OpTy::Adaptor adaptor,
+      ConversionPatternRewriter& rewriter) const final {
+    auto loc = op.getLoc();
+    // Find maximum rank / number of loops.
+    auto getRank = [](Value v) {
+      return v.getType().cast<ShapedType>().getRank();
+    };
+    auto isScalar = [&](Value v) { return getRank(v) == 0; };
+    auto it = llvm::find_if_not(adaptor.getOperands(), isScalar);
+    Value maxRankArg =
+        it != adaptor.getOperands().end() ? *it : adaptor.getOperands().front();
+    int64_t nloops = getRank(maxRankArg);
+
+    // Apply only if all operands are scalar or have the same rank. Some ops,
+    // like `mhlo.select`, support implicit broadcasting of scalars.
+    if (!llvm::all_of(adaptor.getOperands(), [&](Value v) {
+          int64_t r = getRank(v);
+          return r == 0 || r == nloops;
+        })) {
+      return rewriter.notifyMatchFailure(
+          op, "Operands must be os same rank or scalar.");
+    }
+
+    // Find result type, if on tensors.
+    std::optional<ShapedType> resultTy;
+    resultTy = this->typeConverter->convertType(op->getResultTypes().front())
+                   .template dyn_cast<ShapedType>();
+
+    // Check result type compatibility.
+    if (!resultTy || !resultTy->hasRank() || resultTy->getRank() != nloops ||
+        !(resultTy->getElementType().isSignlessIntOrFloat() ||
+          resultTy->getElementType().isa<ComplexType>())) {
+      return rewriter.notifyMatchFailure(
+          op, "mismatched operand/result types or iterator count");
+    }
+
+    if (allOperandsAreScalarTensors(op) && isInBodyOfLinalgOps(op))
+      return failure();
+
+    // Find input/output values and types.
+    ValueRange inputs = adaptor.getOperands();
+    Value output =
+        getEmptyTensorFor(rewriter, loc, *resultTy, op, adaptor.getOperands());
+
+    // Create indexing maps.
+    AffineMap scalarMap = AffineMap::get(nloops, 0, rewriter.getContext());
+    AffineMap idMap = rewriter.getMultiDimIdentityMap(nloops);
+    SmallVector<AffineMap, 4> maps;
+    for (Value v : inputs) maps.push_back(isScalar(v) ? scalarMap : idMap);
+    maps.push_back(idMap);
+
+    // Build `linalg.generic` op.
+    bool failed = false;
+    auto linalgOp = rewriter.create<linalg::GenericOp>(
+        loc, resultTy ? *resultTy : TypeRange{}, inputs, output, maps,
+        getNParallelLoopsAttrs(nloops),
+        [&](OpBuilder& nestedBuilder, Location /*nested_loc*/,
+            ValueRange args) {
+          Type innerResultTy = getElementTypeOrSelf(output);
+          auto argvec = llvm::to_vector<2>(args.take_front(inputs.size()));
+          auto semiring = preSparsify(op, argvec, innerResultTy, &rewriter);
+          Value innerResult = mlir::stablehlo::StableHloOpToStdScalarOp::mapOp(
+              op, innerResultTy, argvec, &rewriter);
+          if (innerResult == nullptr) {
+            failed = true;
+          } else {
+            innerResult = postSparsify(op, semiring, innerResult, &rewriter);
+            nestedBuilder.create<linalg::YieldOp>(loc, innerResult);
+          }
+        },
+        linalg::getPrunedAttributeList(op));
+    if (failed) return failure();
+
+    rewriter.replaceOp(op, linalgOp->getResults());
+    return success();
+  }
+};
+
+}  // namespace mlir::iree_compiler::stablehlo
+
+#endif  // IREE_COMPILER_INPUTCONVERSION_STABLEHLO_LEGALIZE_TO_LINALG_UTILS_H_
diff --git a/compiler/src/iree/compiler/InputConversion/StableHLO/MapStableHLOToScalarOp.h b/compiler/src/iree/compiler/InputConversion/StableHLO/MapStableHLOToScalarOp.h
new file mode 100644
index 000000000000..43e35926bc0c
--- /dev/null
+++ b/compiler/src/iree/compiler/InputConversion/StableHLO/MapStableHLOToScalarOp.h
@@ -0,0 +1,1319 @@
+// Copyright 2019 The IREE Authors
+//
+// Licensed under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+#ifndef IREE_COMPILER_INPUTCONVERSION_STABLEHLO_MAP_STABLEHLO_TO_SCALAR_OP_H
+#define IREE_COMPILER_INPUTCONVERSION_STABLEHLO_MAP_STABLEHLO_TO_SCALAR_OP_H
+
+#include <optional>
+
+#include "llvm/ADT/ArrayRef.h"
+#include "llvm/ADT/StringRef.h"
+#include "llvm/ADT/StringSwitch.h"
+#include "mlir/Dialect/Arith/IR/Arith.h"
+#include "mlir/Dialect/Complex/IR/Complex.h"
+#include "mlir/Dialect/Math/IR/Math.h"
+#include "mlir/Dialect/SCF/IR/SCF.h"
+#include "mlir/Dialect/Vector/IR/VectorOps.h"
+#include "mlir/IR/BuiltinTypes.h"
+#include "mlir/IR/ImplicitLocOpBuilder.h"
+#include "mlir/IR/TypeUtilities.h"
+#include "stablehlo/dialect/StablehloOps.h"
+
+namespace mlir {
+namespace stablehlo {
+namespace impl {
+
+// A struct to map StableHloBinaryOpTy type to the corresponding floating-point
+// and integer scalar operation types.
+template <typename StableHloBinaryOpTy>
+struct StableHloToScalarOp {
+  using FOp = void;
+  using IOp = void;
+  using UOp = void;
+  using COp = void;
+};
+
+template <>
+struct StableHloToScalarOp<stablehlo::AddOp> {
+  using FOp = ::mlir::arith::AddFOp;
+  using IOp = ::mlir::arith::AddIOp;
+  using UOp = ::mlir::arith::AddIOp;
+  using COp = ::mlir::complex::AddOp;
+};
+template <>
+struct StableHloToScalarOp<stablehlo::AndOp> {
+  using IOp = ::mlir::arith::AndIOp;
+  using UOp = ::mlir::arith::AndIOp;
+};
+template <>
+struct StableHloToScalarOp<stablehlo::CbrtOp> {
+  using FOp = ::mlir::math::CbrtOp;
+};
+template <>
+struct StableHloToScalarOp<stablehlo::CompareOp> {
+  using FOp = ::mlir::arith::CmpFOp;
+  using IOp = ::mlir::arith::CmpIOp;
+  using UOp = ::mlir::arith::CmpIOp;
+};
+template <>
+struct StableHloToScalarOp<stablehlo::CeilOp> {
+  using FOp = ::mlir::math::CeilOp;
+};
+template <>
+struct StableHloToScalarOp<stablehlo::ClzOp> {
+  using IOp = ::mlir::math::CountLeadingZerosOp;
+  using UOp = ::mlir::math::CountLeadingZerosOp;
+};
+template <>
+struct StableHloToScalarOp<stablehlo::CosineOp> {
+  using FOp = ::mlir::math::CosOp;
+  using COp = ::mlir::complex::CosOp;
+};
+template <>
+struct StableHloToScalarOp<stablehlo::ExpOp> {
+  using FOp = ::mlir::math::ExpOp;
+  using COp = ::mlir::complex::ExpOp;
+};
+template <>
+struct StableHloToScalarOp<stablehlo::Expm1Op> {
+  using FOp = ::mlir::math::ExpM1Op;
+  using COp = ::mlir::complex::Expm1Op;
+};
+template <>
+struct StableHloToScalarOp<stablehlo::FloorOp> {
+  using FOp = ::mlir::math::FloorOp;
+};
+template <>
+struct StableHloToScalarOp<stablehlo::LogOp> {
+  using FOp = ::mlir::math::LogOp;
+  using COp = ::mlir::complex::LogOp;
+};
+template <>
+struct StableHloToScalarOp<stablehlo::Log1pOp> {
+  using FOp = ::mlir::math::Log1pOp;
+  using COp = ::mlir::complex::Log1pOp;
+};
+template <>
+struct StableHloToScalarOp<stablehlo::MulOp> {
+  using FOp = ::mlir::arith::MulFOp;
+  using IOp = ::mlir::arith::MulIOp;
+  using UOp = ::mlir::arith::MulIOp;
+  using COp = ::mlir::complex::MulOp;
+};
+template <>
+struct StableHloToScalarOp<stablehlo::OrOp> {
+  using IOp = ::mlir::arith::OrIOp;
+  using UOp = ::mlir::arith::OrIOp;
+};
+template <>
+struct StableHloToScalarOp<stablehlo::PopulationCountOp> {
+  using IOp = ::mlir::math::CtPopOp;
+  using UOp = ::mlir::math::CtPopOp;
+};
+template <>
+struct StableHloToScalarOp<stablehlo::RsqrtOp> {
+  using FOp = ::mlir::math::RsqrtOp;
+  using COp = ::mlir::complex::RsqrtOp;
+};
+template <>
+struct StableHloToScalarOp<stablehlo::RoundNearestEvenOp> {
+  using FOp = ::mlir::math::RoundEvenOp;
+};
+template <>
+struct StableHloToScalarOp<stablehlo::RoundOp> {
+  using FOp = ::mlir::math::RoundOp;
+};
+template <>
+struct StableHloToScalarOp<stablehlo::SubtractOp> {
+  using FOp = ::mlir::arith::SubFOp;
+  using IOp = ::mlir::arith::SubIOp;
+  using UOp = ::mlir::arith::SubIOp;
+  using COp = ::mlir::complex::SubOp;
+};
+template <>
+struct StableHloToScalarOp<stablehlo::SqrtOp> {
+  using FOp = ::mlir::math::SqrtOp;
+  using COp = ::mlir::complex::SqrtOp;
+};
+template <>
+struct StableHloToScalarOp<stablehlo::SineOp> {
+  using FOp = ::mlir::math::SinOp;
+  using COp = ::mlir::complex::SinOp;
+};
+// FIXME(Jakub)
+/*
+template <>
+struct StableHloToScalarOp<stablehlo::TanOp> {
+  using FOp = ::mlir::math::TanOp;
+  using COp = ::mlir::complex::TanOp;
+};
+*/
+template <>
+struct StableHloToScalarOp<stablehlo::Atan2Op> {
+  using FOp = ::mlir::math::Atan2Op;
+  using COp = ::mlir::complex::Atan2Op;
+};
+template <>
+struct StableHloToScalarOp<stablehlo::TanhOp> {
+  using FOp = ::mlir::math::TanhOp;
+  using COp = ::mlir::complex::TanhOp;
+};
+template <>
+struct StableHloToScalarOp<stablehlo::XorOp> {
+  using IOp = ::mlir::arith::XOrIOp;
+  using UOp = ::mlir::arith::XOrIOp;
+};
+
+// Alias for the map from StableHLO binary op type to STD floating-point op
+// type.
+template <typename StableHloOp>
+using ScalarFOp = typename StableHloToScalarOp<StableHloOp>::FOp;
+// Alias for the map from StableHLO binary op type to STD signed integer op
+// type.
+template <typename StableHloOp>
+using ScalarIOp = typename StableHloToScalarOp<StableHloOp>::IOp;
+// Alias for the map from StableHLO binary op type to STD unsigned integer op
+// type.
+template <typename StableHloOp>
+using ScalarUOp = typename StableHloToScalarOp<StableHloOp>::UOp;
+// Alias for the map from StableHLO binary op type to STD complex op type.
+template <typename StableHloOp>
+using ScalarCOp = typename StableHloToScalarOp<StableHloOp>::COp;
+
+template <typename... Args>
+struct MapStableHloOpToScalarOpImpl {
+  Value operator()(Location /*loc*/, ArrayRef<Type> /*ResultTypes*/,
+                   ArrayRef<Type> /*argTypes*/, ValueRange /*args*/,
+                   OpBuilder* /*b*/) {
+    return nullptr;
+  }
+};
+
+template <typename StdScalarOp>
+struct MapStableHloOpToScalarOpImpl<StdScalarOp> {
+  Value operator()(Location loc, ArrayRef<Type> resultTypes,
+                   ArrayRef<Type> /*argTypes*/, ValueRange args, OpBuilder* b) {
+    return b->template create<StdScalarOp>(loc, resultTypes, args,
+                                           std::nullopt);
+  }
+};
+
+template <typename SupportedType, typename StdScalarOp, typename... Args>
+struct MapStableHloOpToScalarOpImpl<SupportedType, StdScalarOp, Args...> {
+  Value operator()(Location loc, ArrayRef<Type> resultTypes,
+                   ArrayRef<Type> argTypes, ValueRange args, OpBuilder* b) {
+    Type elementType = getElementTypeOrSelf(argTypes.front());
+    if (SupportedType{}(elementType)) {
+      return b->template create<StdScalarOp>(loc, resultTypes, args,
+                                             std::nullopt);
+    }
+    return MapStableHloOpToScalarOpImpl<Args...>{}(loc, resultTypes, argTypes,
+                                                   args, b);
+  }
+};
+
+template <typename SupportedType, typename... Args>
+struct MapStableHloOpToScalarOpImpl<SupportedType, void, Args...> {
+  Value operator()(Location loc, ArrayRef<Type> resultTypes,
+                   ArrayRef<Type> argTypes, ValueRange args, OpBuilder* b) {
+    return MapStableHloOpToScalarOpImpl<Args...>{}(loc, resultTypes, argTypes,
+                                                   args, b);
+  }
+};
+
+struct IsAnyIntegerType {
+  bool operator()(Type t) { return t.isa<IntegerType>(); }
+};
+
+struct IsSignedIntegerType {
+  bool operator()(Type t) {
+    // Pretend that signless is signed. This will change eventually.
+    return t.isa<IntegerType>() && !t.isUnsignedInteger() &&
+           !t.isSignlessInteger(1);
+  }
+};
+
+struct IsUnsignedIntegerType {
+  bool operator()(Type t) {
+    return t.isUnsignedInteger() || t.isSignlessInteger(1);
+  }
+};
+
+struct IsFloatType {
+  bool operator()(Type t) { return t.isa<FloatType>(); }
+};
+
+struct IsComplexType {
+  bool operator()(Type t) { return t.isa<ComplexType>(); }
+};
+
+template <template <typename T> class MapTy, typename OpTy,
+          typename PredTy = llvm::is_detected<MapTy, OpTy>>
+struct MapableIf {
+  using type = void;
+};
+template <template <typename T> class MapTy, typename OpTy>
+struct MapableIf<MapTy, OpTy, std::true_type> {
+  using type = MapTy<OpTy>;
+};
+
+// Inserts the computation that corresponds to the body of the loop for lowered
+// StableHLO unary/binary op. Returns the value for the result.
+template <typename StableHloOpTy>
+inline Value mapStableHloOpToStdScalarOp(
+    Location loc, ArrayRef<Type> resultTypes, ArrayRef<Type> argTypes,
+    typename StableHloOpTy::Adaptor adaptor, OpBuilder* b) {
+  using ScalarIOpOrVoid = typename MapableIf<ScalarIOp, StableHloOpTy>::type;
+  using ScalarUOpOrVoid = typename MapableIf<ScalarUOp, StableHloOpTy>::type;
+  using ScalarFOpOrVoid = typename MapableIf<ScalarFOp, StableHloOpTy>::type;
+  using ScalarCOpOrVoid = typename MapableIf<ScalarCOp, StableHloOpTy>::type;
+  return MapStableHloOpToScalarOpImpl<IsSignedIntegerType, ScalarIOpOrVoid,
+                                      IsUnsignedIntegerType, ScalarUOpOrVoid,
+                                      IsFloatType, ScalarFOpOrVoid,
+                                      IsComplexType, ScalarCOpOrVoid>{}(
+      loc, resultTypes, argTypes, adaptor.getOperands(), b);
+}
+
+template <>
+inline Value mapStableHloOpToStdScalarOp<stablehlo::AbsOp>(
+    Location loc, ArrayRef<Type> resultTypes, ArrayRef<Type> argTypes,
+    stablehlo::AbsOp::Adaptor adaptor, OpBuilder* b) {
+  Type elementType = getElementTypeOrSelf(argTypes.front());
+  if (elementType.isa<FloatType>()) {
+    return MapStableHloOpToScalarOpImpl<IsFloatType, ::mlir::math::AbsFOp>{}(
+        loc, resultTypes, argTypes, adaptor.getOperands(), b);
+  }
+  if (elementType.isa<ComplexType>()) {
+    return MapStableHloOpToScalarOpImpl<IsComplexType,
+                                        ::mlir::complex::AbsOp>{}(
+        loc, resultTypes, argTypes, adaptor.getOperands(), b);
+  }
+  if (elementType.isSignlessInteger() || elementType.isSignedInteger()) {
+    // lmhlo.abs(x, result) ->  result = select((x > 0), x, sub(0, x))
+    Value lhs = adaptor.getOperand();
+    Value zeroIntval =
+        b->create<arith::ConstantOp>(loc, b->getZeroAttr(lhs.getType()));
+    auto lhsGtZero = b->create<ScalarIOp<CompareOp>>(
+        loc, arith::CmpIPredicate::sge, lhs, zeroIntval);
+    auto negVal =
+        b->create<ScalarIOp<stablehlo::SubtractOp>>(loc, zeroIntval, lhs);
+    return b->create<::mlir::arith::SelectOp>(loc, lhsGtZero, lhs, negVal);
+  }
+  return nullptr;
+}
+
+// Return a constant for v of type t, splat if t is a vector type.
+inline Value getConstantOrSplat(OpBuilder* b, Location loc, Type t,
+                                Attribute v) {
+  if (VectorType vecType = t.dyn_cast<VectorType>()) {
+    v = SplatElementsAttr::get(vecType, v);
+  }
+  return b->create<arith::ConstantOp>(loc, t, v);
+}
+
+template <typename PredicateType>
+inline std::optional<PredicateType> getCmpPredicate(
+    stablehlo::ComparisonDirection, bool) {
+  return std::nullopt;
+}
+
+template <>
+inline std::optional<arith::CmpFPredicate>
+getCmpPredicate<arith::CmpFPredicate>(
+    stablehlo::ComparisonDirection comparisonDirection, bool isSigned) {
+  assert(isSigned && "cannot have an unsigned float!");
+  return llvm::StringSwitch<std::optional<arith::CmpFPredicate>>(
+             stringifyComparisonDirection(comparisonDirection))
+      .Case("EQ", arith::CmpFPredicate::OEQ)
+      .Case("NE", arith::CmpFPredicate::UNE)
+      .Case("GE", arith::CmpFPredicate::OGE)
+      .Case("GT", arith::CmpFPredicate::OGT)
+      .Case("LE", arith::CmpFPredicate::OLE)
+      .Case("LT", arith::CmpFPredicate::OLT)
+      .Default(std::nullopt);
+}
+
+template <>
+inline std::optional<arith::CmpIPredicate>
+getCmpPredicate<arith::CmpIPredicate>(
+    stablehlo::ComparisonDirection comparisonDirection, bool isSigned) {
+  return llvm::StringSwitch<std::optional<arith::CmpIPredicate>>(
+             stringifyComparisonDirection(comparisonDirection))
+      .Case("EQ", arith::CmpIPredicate::eq)
+      .Case("NE", arith::CmpIPredicate::ne)
+      .Case("GE",
+            isSigned ? arith::CmpIPredicate::sge : arith::CmpIPredicate::uge)
+      .Case("GT",
+            isSigned ? arith::CmpIPredicate::sgt : arith::CmpIPredicate::ugt)
+      .Case("LE",
+            isSigned ? arith::CmpIPredicate::sle : arith::CmpIPredicate::ule)
+      .Case("LT",
+            isSigned ? arith::CmpIPredicate::slt : arith::CmpIPredicate::ult)
+      .Default(std::nullopt);
+}
+
+inline Value cmpComplex(Location loc, Value lhs, Value rhs,
+                        stablehlo::ComparisonDirection comparisonDirection,
+                        OpBuilder* b) {
+  auto complexType = lhs.getType().cast<ComplexType>();
+  if (complexType.getElementType().isa<FloatType>()) {
+    if (comparisonDirection == stablehlo::ComparisonDirection::EQ) {
+      return b->create<complex::EqualOp>(loc, lhs, rhs);
+    }
+    if (comparisonDirection == stablehlo::ComparisonDirection::NE) {
+      return b->create<complex::NotEqualOp>(loc, lhs, rhs);
+    }
+
+    // Perform a lexicographical comparison for the (real, imaginary) pair.
+    Type complexFloatTy = complexType.getElementType();
+
+    Value lhsReal = b->create<complex::ReOp>(loc, complexFloatTy, lhs);
+    Value rhsReal = b->create<complex::ReOp>(loc, complexFloatTy, rhs);
+
+    Value lhsImag = b->create<complex::ImOp>(loc, complexFloatTy, lhs);
+    Value rhsImag = b->create<complex::ImOp>(loc, complexFloatTy, rhs);
+
+    auto predicate = getCmpPredicate<arith::CmpFPredicate>(comparisonDirection,
+                                                           /*is_signed=*/true);
+    assert(predicate.has_value() && "expected valid comparison direction");
+
+    //   if (lhsReal == rhsReal && lhsImag `predicate` rhsImag ||
+    //       lhsReal `predicate` rhsReal)
+    Value realsAreEq = b->create<arith::CmpFOp>(loc, arith::CmpFPredicate::OEQ,
+                                                lhsReal, rhsReal);
+    Value imagsAreOrdered =
+        b->create<arith::CmpFOp>(loc, *predicate, lhsImag, rhsImag);
+    Value realsAreOrdered =
+        b->create<arith::CmpFOp>(loc, *predicate, lhsReal, rhsReal);
+
+    Value orLhs = b->create<arith::AndIOp>(loc, realsAreEq, imagsAreOrdered);
+    return b->create<arith::OrIOp>(loc, orLhs, realsAreOrdered);
+  }
+  return nullptr;
+}
+
+template <>
+inline Value mapStableHloOpToStdScalarOp<stablehlo::CompareOp>(
+    Location loc, ArrayRef<Type> /*resultTypes*/, ArrayRef<Type> argTypes,
+    stablehlo::CompareOp::Adaptor adaptor, OpBuilder* b) {
+  stablehlo::ComparisonDirection comparisonDirection =
+      adaptor.getComparisonDirection();
+  const auto& lhs = adaptor.getLhs();
+  const auto& rhs = adaptor.getRhs();
+  Type elementType = getElementTypeOrSelf(argTypes.front());
+  if (elementType.isa<IntegerType>()) {
+    bool isUnsigned = IsUnsignedIntegerType{}(elementType);
+    std::optional<arith::CmpIPredicate> predicate =
+        getCmpPredicate<arith::CmpIPredicate>(comparisonDirection, !isUnsigned);
+    assert(predicate.has_value() && "expected valid comparison direction");
+    return b->create<ScalarIOp<stablehlo::CompareOp>>(loc, predicate.value(),
+                                                      lhs, rhs);
+  }
+  if (auto floatType = elementType.dyn_cast<FloatType>()) {
+    if (adaptor.getCompareType() &&
+        *adaptor.getCompareType() == stablehlo::ComparisonType::TOTALORDER) {
+      // The semantics of totalorder fp compare are
+      // -NaN < -Inf < -Finite < -0 < +0 < +Finite < +Inf < +NaN
+      auto intType = b->getIntegerType(floatType.getWidth());
+      auto zero =
+          b->create<arith::ConstantOp>(loc, intType, b->getZeroAttr(intType));
+      auto max = b->create<arith::ConstantOp>(
+          loc, intType,
+          b->getIntegerAttr(intType,
+                            APInt::getSignedMaxValue(floatType.getWidth())));
+      // Switch from a floating point value to a integer value in such a way
+      // that when using the integer value to compare, we get the same result
+      // for normal values, and -NaN is treated as the smallest value, and NaN
+      // is treated as the largest value.
+      // If f is a float, and
+      // x = bit_cast<int32_t>(f);
+      // y = x < 0 ? numeric_limits<int32_t>::max() - x : x;
+      // then y is ordered as an int32_t such that finite values have the
+      // obvious order, -0 is ordered before 0, and -NaN and NaN appear at the
+      // beginning and end of the ordering.
+      auto toIntegral = [&](Value v) {
+        auto x = b->create<arith::BitcastOp>(loc, intType, v);
+        auto cmp =
+            b->create<arith::CmpIOp>(loc, arith::CmpIPredicate::slt, x, zero);
+        auto sub = b->create<arith::SubIOp>(loc, max, x);
+        return b->create<arith::SelectOp>(loc, cmp, sub, x);
+      };
+      auto lhsInt = toIntegral(lhs);
+      auto rhsInt = toIntegral(rhs);
+      auto predicate =
+          getCmpPredicate<arith::CmpIPredicate>(comparisonDirection,
+                                                /*is_signed=*/true);
+      assert(predicate.has_value() && "expected valid comparison direction");
+      return b->create<arith::CmpIOp>(loc, *predicate, lhsInt, rhsInt);
+    }
+    std::optional<arith::CmpFPredicate> predicate =
+        getCmpPredicate<arith::CmpFPredicate>(comparisonDirection,
+                                              /*is_signed=*/true);
+    assert(predicate.has_value() && "expected valid comparison direction");
+    return b->create<ScalarFOp<stablehlo::CompareOp>>(loc, predicate.value(),
+                                                      lhs, rhs);
+  }
+  if (auto complexType = elementType.dyn_cast<ComplexType>())
+    return cmpComplex(loc, lhs, rhs, comparisonDirection, b);
+  return nullptr;
+}
+
+template <>
+inline Value mapStableHloOpToStdScalarOp<stablehlo::ReducePrecisionOp>(
+    Location loc, ArrayRef<Type> /*resultTypes*/, ArrayRef<Type> argTypes,
+    stablehlo::ReducePrecisionOp::Adaptor adaptor, OpBuilder* builder) {
+  using llvm::APInt;
+  mlir::ImplicitLocOpBuilder b(loc, *builder);
+
+  // Integer and float types for casting and constant generation.
+  auto floatType =
+      argTypes.front().cast<TensorType>().getElementType().cast<FloatType>();
+  int64_t nbits = floatType.getWidth();
+  auto intType = mlir::IntegerType::get(loc.getContext(), floatType.getWidth());
+
+  Value xAsInt = b.create<arith::BitcastOp>(intType, adaptor.getOperand());
+
+  // SignificandWidth includes the implicit extra bit.
+  auto srcMantissaBits = floatType.getFPMantissaWidth() - 1;
+  int srcExponentBits = nbits - 1 - srcMantissaBits;
+
+  // Clear the sign bit, it does not participate in rounding and we will restore
+  // it later.
+  APInt signBitMask(nbits, 1);
+  signBitMask <<= nbits - 1;
+
+  APInt expBitsMask(nbits, 1);
+  expBitsMask = ((expBitsMask << srcExponentBits) - 1) << srcMantissaBits;
+
+  auto createConstant = [&](const APInt& v) {
+    return b.create<arith::ConstantIntOp>(v.getZExtValue(), intType)
+        .getResult();
+  };
+
+  Value xAbsBits =
+      b.create<arith::AndIOp>(xAsInt, createConstant(~signBitMask));
+  Value xIsNan = b.create<arith::CmpIOp>(arith::CmpIPredicate::ugt, xAbsBits,
+                                         createConstant(expBitsMask));
+
+  int destMantissaBits = adaptor.getMantissaBits();
+  if (destMantissaBits < static_cast<int>(srcMantissaBits)) {
+    // Last remaining mantissa bit.
+    APInt lastMantissaBitMask(nbits, 1);
+    lastMantissaBitMask <<= srcMantissaBits - destMantissaBits;
+
+    // Compute rounding bias for round-to-nearest with ties to even.  This is
+    // equal to a base value of 0111... plus one bit if the last remaining
+    // mantissa bit is 1.
+    APInt baseRoundingBias = lastMantissaBitMask.lshr(1) - 1;
+
+    Value mantissaDiff = b.create<arith::ConstantIntOp>(
+        srcMantissaBits - destMantissaBits, intType);
+    Value highestMantissaMaskVal = createConstant(lastMantissaBitMask);
+    Value baseRoundingBiasVal = createConstant(baseRoundingBias);
+    Value xLastMantissaBit = b.create<arith::ShRUIOp>(
+        b.create<arith::AndIOp>(xAsInt, highestMantissaMaskVal), mantissaDiff);
+    Value xRoundingBias =
+        b.create<arith::AddIOp>(xLastMantissaBit, baseRoundingBiasVal);
+
+    // Add rounding bias, and mask out truncated bits.  Note that the case
+    // where adding the rounding bias overflows into the exponent bits is
+    // correct; the non-masked mantissa bits will all be zero, and the
+    // exponent will be incremented by one.
+    APInt truncationMask = ~(lastMantissaBitMask - 1);
+    Value xRounded = b.create<arith::AddIOp>(xAsInt, xRoundingBias);
+    xAsInt = b.create<arith::AndIOp>(xRounded, createConstant(truncationMask));
+  }
+
+  int destExponentBits = adaptor.getExponentBits();
+  if (destExponentBits < srcExponentBits) {
+    // An exponent of 2^(n-1)-1 -- that is, 0111... with the zero in the most-
+    // significant bit -- is equal to 1.0f for all exponent sizes.  Adding
+    // 2^(n-1)-1 to this gives us the highest non-infinite exponent for a bit-
+    // size of n, and subtracting 2^(n-1)-1 from this gives us the lowest'
+    // exponent (corresponding to 0.0f).
+    //
+    // Thus, the f32 exponent corresponding to the highest non-infinite
+    // exponent for a bit size of n is (2^7-1) + 2^(n-1)-1, and the f32
+    // exponent corresponding to the lowest exponent for a bit size of n is
+    // (2^7-1) - 2^(n-1)-1.
+    //
+    // Note that we have already checked that exponents_bits >= 1.
+    APInt exponentBias(nbits, 1);
+    exponentBias = (exponentBias << (srcExponentBits - 1)) - 1;
+
+    APInt reducedExponentBias(nbits, 1);
+    reducedExponentBias = (reducedExponentBias << (destExponentBits - 1)) - 1;
+
+    APInt reducedMaxExponent = exponentBias + reducedExponentBias;
+    APInt reducedMinExponent = exponentBias - reducedExponentBias;
+
+    // Do we overflow or underflow?
+    Value xExponent =
+        b.create<arith::AndIOp>(xAsInt, createConstant(expBitsMask));
+    Value xOverflows = b.create<arith::CmpIOp>(
+        arith::CmpIPredicate::ugt, xExponent,
+        createConstant(reducedMaxExponent << srcMantissaBits));
+    Value xUnderflows = b.create<arith::CmpIOp>(
+        arith::CmpIPredicate::ule, xExponent,
+        createConstant(reducedMinExponent << srcMantissaBits));
+
+    // Compute appropriately-signed values of zero and infinity.
+    Value xSignedZero =
+        b.create<arith::AndIOp>(xAsInt, createConstant(signBitMask));
+    Value xSignedInf =
+        b.create<arith::OrIOp>(xSignedZero, createConstant(expBitsMask));
+
+    // Force to zero or infinity if overflow or underflow.  (Note that this
+    // truncates all denormal values to zero, rather than rounding them.)
+    xAsInt = b.create<arith::SelectOp>(xOverflows, xSignedInf, xAsInt);
+    xAsInt = b.create<arith::SelectOp>(xUnderflows, xSignedZero, xAsInt);
+  }
+
+  Value result = b.create<arith::BitcastOp>(floatType, xAsInt);
+  return b.create<arith::SelectOp>(xIsNan, adaptor.getOperand(), result);
+}
+
+// FIXME(Jakub)
+// template <>
+// inline Value mapStableHloOpToStdScalarOp<stablehlo::CopyOp>(
+//     Location /*loc*/, ArrayRef<Type> /*ResultTypes*/,
+//     ArrayRef<Type> /*argTypes*/, stablehlo::CopyOp::Adaptor adaptor,
+//     OpBuilder* /*b*/) {
+//   return adaptor.getOperands().front();
+// }
+
+template <>
+inline Value mapStableHloOpToStdScalarOp<stablehlo::ComplexOp>(
+    Location loc, ArrayRef<Type> resultTypes, ArrayRef<Type> argTypes,
+    stablehlo::ComplexOp::Adaptor adaptor, OpBuilder* b) {
+  return MapStableHloOpToScalarOpImpl<complex::CreateOp>{}(
+      loc, resultTypes, argTypes, adaptor.getOperands(), b);
+}
+
+template <>
+inline Value mapStableHloOpToStdScalarOp<stablehlo::MaxOp>(
+    Location loc, ArrayRef<Type> resultTypes, ArrayRef<Type> argTypes,
+    stablehlo::MaxOp::Adaptor adaptor, OpBuilder* b) {
+  ValueRange operands = adaptor.getOperands();
+  Value lhs = operands.front();
+  Type complexTy = lhs.getType();
+
+  if (!complexTy.isa<ComplexType>())
+    return MapStableHloOpToScalarOpImpl<
+        IsFloatType, arith::MaxFOp, IsSignedIntegerType, arith::MaxSIOp,
+        IsUnsignedIntegerType, arith::MaxUIOp>{}(loc, resultTypes, argTypes,
+                                                 adaptor.getOperands(), b);
+
+  assert(resultTypes.size() == 1 && "MaxOp should return a single result");
+  assert(operands.size() == 2 && "MaxOp should take exactly two arguments");
+
+  Value rhs = operands.back();
+  // 'max' performs a lexicographical comparison for the (real, imaginary) pair.
+  Value cond = cmpComplex(loc, lhs, rhs, stablehlo::ComparisonDirection::GE, b);
+
+  return b->create<arith::SelectOp>(loc, cond, lhs, rhs).getResult();
+}
+
+template <>
+inline Value mapStableHloOpToStdScalarOp<stablehlo::MinOp>(
+    Location loc, ArrayRef<Type> resultTypes, ArrayRef<Type> argTypes,
+    stablehlo::MinOp::Adaptor adaptor, OpBuilder* b) {
+  ValueRange operands = adaptor.getOperands();
+  Value lhs = operands.front();
+  Type complexTy = lhs.getType();
+
+  if (!complexTy.isa<ComplexType>())
+    return MapStableHloOpToScalarOpImpl<
+        IsFloatType, arith::MinFOp, IsSignedIntegerType, arith::MinSIOp,
+        IsUnsignedIntegerType, arith::MinUIOp>{}(loc, resultTypes, argTypes,
+                                                 adaptor.getOperands(), b);
+
+  assert(resultTypes.size() == 1 && "MinOp should return a single result");
+  assert(operands.size() == 2 && "MinOp should take exactly two arguments");
+
+  Value rhs = operands.back();
+  // 'min' performs a lexicographical comparison for the (real, imaginary) pair.
+  Value cond = cmpComplex(loc, lhs, rhs, stablehlo::ComparisonDirection::LE, b);
+
+  return b->create<arith::SelectOp>(loc, cond, lhs, rhs).getResult();
+}
+
+template <>
+inline Value mapStableHloOpToStdScalarOp<stablehlo::RealOp>(
+    Location loc, ArrayRef<Type> resultTypes, ArrayRef<Type> argTypes,
+    stablehlo::RealOp::Adaptor adaptor, OpBuilder* b) {
+  if (!adaptor.getOperand().getType().isa<ComplexType>())
+    return adaptor.getOperand();
+  return MapStableHloOpToScalarOpImpl<complex::ReOp>{}(
+      loc, resultTypes, argTypes, adaptor.getOperands(), b);
+}
+
+template <>
+inline Value mapStableHloOpToStdScalarOp<stablehlo::ImagOp>(
+    Location loc, ArrayRef<Type> resultTypes, ArrayRef<Type> argTypes,
+    stablehlo::ImagOp::Adaptor adaptor, OpBuilder* b) {
+  if (!adaptor.getOperand().getType().isa<ComplexType>())
+    return b->create<arith::ConstantOp>(
+        loc, b->getZeroAttr(adaptor.getOperand().getType()));
+  return MapStableHloOpToScalarOpImpl<complex::ImOp>{}(
+      loc, resultTypes, argTypes, adaptor.getOperands(), b);
+}
+
+// 'target_types' is the unconverted type (signed or unsigned if integer),
+// 'ResultTypes' is the converted type (signless if integer).
+inline Value mapConvertOpToStdScalarOp(Location loc, ArrayRef<Type> targetTypes,
+                                       ArrayRef<Type> resultTypes,
+                                       ArrayRef<Type> argTypes, ValueRange args,
+                                       OpBuilder* b) {
+  assert(targetTypes.size() == 1 && "ConvertOp should return a single result");
+  assert(resultTypes.size() == 1 && "ConvertOp should return a single result");
+  assert(argTypes.size() == 1 && "ConvertOp should take a single argument");
+  assert(args.size() == 1 && "ConvertOp should take a single argument");
+
+  Type sourceType = getElementTypeOrSelf(argTypes.front());
+  Type targetType = getElementTypeOrSelf(targetTypes.front());
+  Type convertedSourceType = getElementTypeOrSelf(args.front());
+
+  // A boolean value is considered to be unsigned when converting to
+  // floating-point. Otherwise, it will become `-1`.
+  if (IsUnsignedIntegerType{}(sourceType) &&
+      mlir::arith::UIToFPOp::areCastCompatible(convertedSourceType,
+                                               targetType)) {
+    return b->create<mlir::arith::UIToFPOp>(loc, resultTypes, args,
+                                            std::nullopt);
+  }
+  if (mlir::arith::SIToFPOp::areCastCompatible(sourceType, targetType)) {
+    return b->create<mlir::arith::SIToFPOp>(loc, resultTypes, args,
+                                            std::nullopt);
+  }
+  if (sourceType.isa<FloatType>() && targetType.isa<FloatType>()) {
+    auto src = sourceType.cast<FloatType>();
+    auto res = targetType.cast<FloatType>();
+    if (src.getWidth() > res.getWidth()) {
+      return b->create<mlir::arith::TruncFOp>(loc, resultTypes, args,
+                                              std::nullopt);
+    }
+    if (src.getWidth() < res.getWidth()) {
+      return b->create<mlir::arith::ExtFOp>(loc, resultTypes, args,
+                                            std::nullopt);
+    }
+    // There's no direct conversion between different 16 bit floating point
+    // types, so go through 32 bit float.
+    if (sourceType != targetType) {
+      assert(sourceType.isBF16() || targetType.isBF16());
+      Value ext = b->create<arith::ExtFOp>(loc, b->getF32Type(), args);
+      return b->create<arith::TruncFOp>(loc, resultTypes, ext);
+    }
+    // No conversion is needed for identical float types.
+    return args.front();
+  }
+  if (targetType.isInteger(/*width=*/1)) {
+    // When casting to bool, we need to compare whether the value is equal to
+    // zero.
+    if (sourceType.isSignlessInteger() || sourceType.isUnsignedInteger()) {
+      Value zeroIntval = b->create<arith::ConstantOp>(
+          loc, b->getZeroAttr(args.front().getType()));
+      return b->create<mlir::arith::CmpIOp>(loc, arith::CmpIPredicate::ne,
+                                            args.front(), zeroIntval);
+    }
+    if (sourceType.isa<FloatType>()) {
+      Value zero = b->create<arith::ConstantOp>(
+          loc, b->getZeroAttr(args.front().getType()));
+      return b->create<mlir::arith::CmpFOp>(loc, arith::CmpFPredicate::UNE,
+                                            args.front(), zero);
+    }
+  }
+  if (sourceType.isa<IntegerType>() && targetType.isa<IntegerType>()) {
+    auto src = sourceType.cast<IntegerType>();
+    auto res = targetType.cast<IntegerType>();
+    if (src.getWidth() > res.getWidth()) {
+      return b->create<mlir::arith::TruncIOp>(loc, resultTypes, args,
+                                              std::nullopt);
+    }
+    if (src.getWidth() < res.getWidth()) {
+      // Special case boolean values, so they get casted to `1` instead of `-1`.
+      if (IsUnsignedIntegerType{}(src)) {
+        return b->create<mlir::arith::ExtUIOp>(loc, resultTypes, args,
+                                               std::nullopt);
+      }
+      return b->create<mlir::arith::ExtSIOp>(loc, resultTypes, args,
+                                             std::nullopt);
+    }
+    // No conversion is needed for the same width integers
+    return args.front();
+  }
+  if (targetType.isUnsignedInteger() &&
+      mlir::arith::FPToUIOp::areCastCompatible(convertedSourceType,
+                                               targetType)) {
+    return b->create<mlir::arith::FPToUIOp>(loc, resultTypes, args,
+                                            std::nullopt);
+  }
+  if (mlir::arith::FPToSIOp::areCastCompatible(convertedSourceType,
+                                               targetType)) {
+    return b->create<mlir::arith::FPToSIOp>(loc, resultTypes, args,
+                                            std::nullopt);
+  }
+  if (targetType.isa<ComplexType>()) {
+    Type targetElementType = targetType.cast<ComplexType>().getElementType();
+    assert(!targetElementType.isa<ComplexType>() &&
+           "elements of complex numbers should not be complex");
+    Value targetReal;
+    Value targetImag;
+    if (sourceType.isa<ComplexType>()) {
+      // We are converting from complex type: convert real and imaginary parts
+      // separately.
+      Type sourceElementType = sourceType.cast<ComplexType>().getElementType();
+      assert(!sourceElementType.isa<ComplexType>() &&
+             "elements of complex numbers should not be complex");
+      Value sourceReal =
+          b->create<mlir::complex::ReOp>(loc, sourceElementType, args.front());
+      targetReal =
+          mapConvertOpToStdScalarOp(loc, targetElementType, targetElementType,
+                                    sourceElementType, sourceReal, b);
+      Value sourceImag =
+          b->create<mlir::complex::ImOp>(loc, sourceElementType, args.front());
+      targetImag =
+          mapConvertOpToStdScalarOp(loc, targetElementType, targetElementType,
+                                    sourceElementType, sourceImag, b);
+    } else {
+      // We are converting from real (float, integer, etc.) type, convert the
+      // real part and set the imaginary part to 0.
+      targetReal = mapConvertOpToStdScalarOp(
+          loc, targetElementType, targetElementType, argTypes, args, b);
+      targetImag = b->create<mlir::arith::ConstantOp>(
+          loc, b->getFloatAttr(targetElementType, 0.0));
+    }
+    return b->create<mlir::complex::CreateOp>(loc, targetType, targetReal,
+                                              targetImag);
+  }
+  if (auto sourceComplexType = sourceType.dyn_cast<ComplexType>()) {
+    auto sourceElementType = sourceComplexType.getElementType();
+    // When converting from complex to a non-complex type, we take just the real
+    // part of the complex number.
+    Value sourceReal =
+        b->create<mlir::complex::ReOp>(loc, sourceElementType, args.front());
+    return mapConvertOpToStdScalarOp(loc, targetTypes, resultTypes,
+                                     sourceElementType, sourceReal, b);
+  }
+  return nullptr;
+}
+
+/// Lower bitcast operations where the input and resulting type are the same
+/// bitwidth, thus implying that the operation is fully defined by parallel
+/// loops and scalar operations without any shape dimension changes.
+template <>
+inline Value mapStableHloOpToStdScalarOp<stablehlo::BitcastConvertOp>(
+    Location loc, ArrayRef<Type> resultTypes, ArrayRef<Type> argTypes,
+    stablehlo::BitcastConvertOp::Adaptor adaptor, OpBuilder* b) {
+  Type argType = getElementTypeOrSelf(argTypes.front());
+  Type resultType = getElementTypeOrSelf(resultTypes.front());
+
+  if (resultType.getIntOrFloatBitWidth() != argType.getIntOrFloatBitWidth())
+    return nullptr;
+
+  return b->create<mlir::arith::BitcastOp>(loc, resultTypes,
+                                           adaptor.getOperands());
+}
+
+template <>
+inline Value mapStableHloOpToStdScalarOp<stablehlo::DotOp>(
+    Location loc, ArrayRef<Type> resultTypes, ArrayRef<Type> argTypes,
+    stablehlo::DotOp::Adaptor adaptor, OpBuilder* b) {
+  // Dot Op converter from lhlo to affine only accepts float and integer types.
+  const auto& lhs = adaptor.getOperands()[0];
+  const auto& rhs = adaptor.getOperands()[1];
+  const auto& result = adaptor.getOperands()[2];
+  Type elementType = lhs.getType();
+  if (elementType.isa<FloatType>()) {
+    Value floatMul =
+        MapStableHloOpToScalarOpImpl<IsFloatType, ::mlir::arith::MulFOp>{}(
+            loc, resultTypes, argTypes, {lhs, rhs}, b);
+    return MapStableHloOpToScalarOpImpl<IsFloatType, ::mlir::arith::AddFOp>{}(
+        loc, resultTypes, argTypes, {floatMul, result}, b);
+  }
+  if (elementType.isa<IntegerType>()) {
+    Value intMul =
+        MapStableHloOpToScalarOpImpl<IsAnyIntegerType, ::mlir::arith::MulIOp>{}(
+            loc, resultTypes, argTypes, {lhs, rhs}, b);
+    return MapStableHloOpToScalarOpImpl<IsAnyIntegerType,
+                                        ::mlir::arith::AddIOp>{}(
+        loc, resultTypes, argTypes, {intMul, result}, b);
+  }
+  return nullptr;
+}
+
+template <>
+inline Value mapStableHloOpToStdScalarOp<stablehlo::IsFiniteOp>(
+    Location loc, ArrayRef<Type> /*ResultTypes*/, ArrayRef<Type> /*argTypes*/,
+    stablehlo::IsFiniteOp::Adaptor adaptor, OpBuilder* b) {
+  if (adaptor.getX().getType().isa<FloatType>()) {
+    auto posInf = APFloat::getInf(
+        adaptor.getX().getType().cast<FloatType>().getFloatSemantics());
+    auto constPosInf = b->create<arith::ConstantOp>(
+        loc, b->getFloatAttr(adaptor.getX().getType(), posInf));
+    Value absX = b->create<::mlir::math::AbsFOp>(loc, adaptor.getX());
+    return b->create<::mlir::arith::CmpFOp>(loc, arith::CmpFPredicate::ONE,
+                                            absX, constPosInf);
+  }
+  return nullptr;
+}
+
+/// Implements the conversion of HLO op to scalar op (to use within region of a
+/// linalg.generic op) for compare-select style operations like min/max.
+template <typename... Args>
+struct CompareSelectOpToStdScalarOp {
+  static Value map(Location /*loc*/,
+                   stablehlo::ComparisonDirection /*comparison_direction*/,
+                   ArrayRef<Type> /*ResultTypes*/, ArrayRef<Type> /*argTypes*/,
+                   ValueRange /*args*/, OpBuilder* /*b*/) {
+    return nullptr;
+  }
+};
+
+/// Specialization which allows converting to a comparison operation in standard
+/// dialect with a given predicate based on the element type of the operand.
+template <typename SupportedType, typename StdCompareOp, typename Predicate,
+          typename... Args>
+struct CompareSelectOpToStdScalarOp<SupportedType, StdCompareOp, Predicate,
+                                    Args...> {
+  static Value map(Location loc,
+                   stablehlo::ComparisonDirection comparisonDirection,
+                   ArrayRef<Type> resultTypes, ArrayRef<Type> argTypes,
+                   ValueRange args, OpBuilder* b) {
+    Type elementType = getElementTypeOrSelf(argTypes.front());
+    if (elementType.isa<SupportedType>()) {
+      auto predicate = getCmpPredicate<Predicate>(
+          comparisonDirection, !elementType.isUnsignedInteger());
+      assert(predicate.has_value() && "expected valid comparison direction");
+      auto cmp = b->template create<StdCompareOp>(loc, predicate.getValue(),
+                                                  args[0], args[1]);
+      return b->create<::mlir::arith::SelectOp>(loc, cmp, args[0], args[1]);
+    }
+    return CompareSelectOpToStdScalarOp<Args...>::map(
+        loc, comparisonDirection, resultTypes, argTypes, args, b);
+  }
+};
+
+inline Value mhloAlwaysPropagateNaN(Value v, ValueRange args, Location loc,
+                                    OpBuilder* b) {
+  Type elementType = getElementTypeOrSelf(args.front().getType());
+  if (auto floatType = elementType.dyn_cast<FloatType>()) {
+    Value isnan = b->create<mlir::arith::CmpFOp>(loc, arith::CmpFPredicate::UNO,
+                                                 args[0], args[1]);
+
+    auto nanApfloat = APFloat::getQNaN(floatType.getFloatSemantics());
+    Value nan = getConstantOrSplat(b, loc, args[0].getType(),
+                                   b->getFloatAttr(floatType, nanApfloat));
+    v = b->create<mlir::arith::SelectOp>(loc, isnan, nan, v);
+  }
+  return v;
+}
+
+template <>
+inline Value mapStableHloOpToStdScalarOp<stablehlo::ClampOp>(
+    Location loc, ArrayRef<Type> resultTypes, ArrayRef<Type> argTypes,
+    stablehlo::ClampOp::Adaptor op, OpBuilder* b) {
+  // clamp(lb, x, ub) = min(max(lb, x), ub)
+  Value maxLbX = mapStableHloOpToStdScalarOp<stablehlo::MaxOp>(
+      loc, resultTypes, argTypes, ValueRange{op.getMin(), op.getOperand()}, b);
+  return mapStableHloOpToStdScalarOp<stablehlo::MinOp>(
+      loc, resultTypes, argTypes, ValueRange{maxLbX, op.getMax()}, b);
+}
+
+template <typename U, typename S>
+inline Value makeSafeIntDiv(ImplicitLocOpBuilder& lb, Type originalType,
+                            Value lhs, Value rhs, Value returnedOnZero,
+                            Value returnedOnSignedOverflow) {
+  Type type = lhs.getType();
+  auto elementType = getElementTypeOrSelf(type).cast<IntegerType>();
+  Value zero = lb.create<arith::ConstantOp>(lb.getZeroAttr(type));
+  auto makeConstant = [&](const APInt& i) {
+    return getConstantOrSplat(&lb, lb.getLoc(), type,
+                              lb.getIntegerAttr(elementType, i));
+  };
+  Value one = makeConstant(APInt(elementType.getWidth(), 1));
+  Value rhsIsZero =
+      lb.create<arith::CmpIOp>(arith::CmpIPredicate::eq, rhs, zero);
+
+  // For unsigned just set the divisor to 1 when it would be 0.
+  if (originalType.isUnsignedInteger()) {
+    Value safeRhs = lb.create<arith::SelectOp>(rhsIsZero, one, rhs);
+    Value safeDiv = lb.create<U>(lhs, safeRhs);
+    return lb.create<arith::SelectOp>(rhsIsZero, returnedOnZero, safeDiv);
+  }
+
+  // For signed also check for INT_MIN / -1.
+  Value smin = makeConstant(APInt::getSignedMinValue(elementType.getWidth()));
+  Value lhsIsSmin =
+      lb.create<arith::CmpIOp>(arith::CmpIPredicate::eq, lhs, smin);
+  Value minusOne = makeConstant(APInt::getAllOnes(elementType.getWidth()));
+  Value rhsIsMinusOne =
+      lb.create<arith::CmpIOp>(arith::CmpIPredicate::eq, rhs, minusOne);
+  Value hasIntMinOverflow = lb.create<arith::AndIOp>(lhsIsSmin, rhsIsMinusOne);
+  Value rhsIsUnsafe = lb.create<arith::OrIOp>(rhsIsZero, hasIntMinOverflow);
+  Value safeRhs = lb.create<arith::SelectOp>(rhsIsUnsafe, one, rhs);
+  Value safeDiv = lb.create<S>(lhs, safeRhs);
+  Value safeSmin = lb.create<arith::SelectOp>(
+      hasIntMinOverflow, returnedOnSignedOverflow, safeDiv);
+  return lb.create<arith::SelectOp>(rhsIsZero, returnedOnZero, safeSmin);
+}
+
+template <>
+inline Value mapStableHloOpToStdScalarOp<stablehlo::DivOp>(
+    Location loc, ArrayRef<Type> resultTypes, ArrayRef<Type> argTypes,
+    stablehlo::DivOp::Adaptor adaptor, OpBuilder* b) {
+  Type originalType = getElementTypeOrSelf(argTypes.front());
+  if (originalType.isa<ComplexType, FloatType>()) {
+    return MapStableHloOpToScalarOpImpl<IsFloatType, arith::DivFOp,
+                                        IsComplexType, complex::DivOp>{}(
+        loc, resultTypes, argTypes, adaptor.getOperands(), b);
+  }
+
+  // Integer division overflow behavior:
+  //
+  // X / 0 == -1
+  // INT_SMIN /s -1 = INT_SMIN
+  ImplicitLocOpBuilder lb(loc, *b);
+  Type type = adaptor.getLhs().getType();
+  auto elementType = getElementTypeOrSelf(type).cast<IntegerType>();
+  auto makeConstant = [&](const APInt& i) {
+    return getConstantOrSplat(&lb, lb.getLoc(), type,
+                              lb.getIntegerAttr(elementType, i));
+  };
+  Value minusOne = makeConstant(APInt::getAllOnes(elementType.getWidth()));
+  Value smin = makeConstant(APInt::getSignedMinValue(elementType.getWidth()));
+  return makeSafeIntDiv<arith::DivUIOp, arith::DivSIOp>(
+      lb, originalType, adaptor.getLhs(), adaptor.getRhs(),
+      /*returnedOnZero=*/minusOne,
+      /*returnedOnSignedOverflow=*/smin);
+}
+
+template <>
+inline Value mapStableHloOpToStdScalarOp<stablehlo::RemOp>(
+    Location loc, ArrayRef<Type> resultTypes, ArrayRef<Type> argTypes,
+    stablehlo::RemOp::Adaptor adaptor, OpBuilder* b) {
+  Type originalType = getElementTypeOrSelf(argTypes.front());
+  if (originalType.isa<ComplexType, FloatType>()) {
+    return MapStableHloOpToScalarOpImpl<IsFloatType, arith::RemFOp>{}(
+        loc, resultTypes, argTypes, adaptor.getOperands(), b);
+  }
+
+  // Integer remainder overflow behavior:
+  //
+  // X % 0 == X
+  // INT_SMIN %s -1 = 0
+  ImplicitLocOpBuilder lb(loc, *b);
+  Type type = adaptor.getLhs().getType();
+  Value zero = lb.create<arith::ConstantOp>(lb.getZeroAttr(type));
+  return makeSafeIntDiv<arith::RemUIOp, arith::RemSIOp>(
+      lb, originalType, adaptor.getLhs(), adaptor.getRhs(),
+      /*returnedOnZero=*/adaptor.getLhs(),
+      /*returnedOnSignedOverflow=*/zero);
+}
+
+template <>
+inline Value mapStableHloOpToStdScalarOp<stablehlo::NegOp>(
+    Location loc, ArrayRef<Type> resultTypes, ArrayRef<Type> argTypes,
+    stablehlo::NegOp::Adaptor adaptor, OpBuilder* b) {
+  Type elementType = getElementTypeOrSelf(adaptor.getOperand().getType());
+  if (elementType.isa<ComplexType, FloatType>()) {
+    return MapStableHloOpToScalarOpImpl<IsFloatType, ::mlir::arith::NegFOp,
+                                        IsComplexType,
+                                        ::mlir::complex::NegOp>{}(
+        loc, resultTypes, argTypes, adaptor.getOperands(), b);
+  }
+  if (elementType.isa<IntegerType>()) {
+    // lmhlo.neg(x, result) -> result = sub(0, x)
+    Value lhs = adaptor.getOperand();
+    Value zeroIntval =
+        b->create<arith::ConstantOp>(loc, b->getZeroAttr(lhs.getType()));
+    return b->create<ScalarIOp<stablehlo::SubtractOp>>(loc, zeroIntval, lhs);
+  }
+  return nullptr;
+}
+
+template <>
+inline Value mapStableHloOpToStdScalarOp<stablehlo::NotOp>(
+    Location loc, ArrayRef<Type> /*ResultTypes*/, ArrayRef<Type> /*argTypes*/,
+    stablehlo::NotOp::Adaptor adaptor, OpBuilder* b) {
+  Type elementType = getElementTypeOrSelf(adaptor.getOperand().getType());
+  if (auto integerType = elementType.dyn_cast<IntegerType>()) {
+    // lmhlo.not(x) -> x ^ -1
+    Value allOnes = getConstantOrSplat(
+        b, loc, adaptor.getOperand().getType(),
+        b->getIntegerAttr(integerType,
+                          APInt::getAllOnes(integerType.getWidth())));
+    return b->create<::mlir::arith::XOrIOp>(loc, allOnes, adaptor.getOperand());
+  }
+  return nullptr;
+}
+
+template <>
+inline Value mapStableHloOpToStdScalarOp<stablehlo::LogisticOp>(
+    Location loc, ArrayRef<Type> resultTypes, ArrayRef<Type> /*argTypes*/,
+    stablehlo::LogisticOp::Adaptor adaptor, OpBuilder* b) {
+  // 1.0 / (1.0 - exp(-x))
+  Value negX = mapStableHloOpToStdScalarOp<stablehlo::NegOp>(
+      loc, resultTypes, resultTypes, {adaptor.getOperand()}, b);
+  Value expNegX = mapStableHloOpToStdScalarOp<stablehlo::ExpOp>(
+      loc, resultTypes, resultTypes, {{negX}}, b);
+
+  Value oneFloat = b->create<arith::ConstantOp>(loc, b->getF32FloatAttr(1.0));
+  Value one = mapConvertOpToStdScalarOp(loc, resultTypes, resultTypes,
+                                        {oneFloat.getType()}, {{oneFloat}}, b);
+  Value oneAddExprNegX = mapStableHloOpToStdScalarOp<stablehlo::AddOp>(
+      loc, resultTypes, resultTypes, {{expNegX, one}}, b);
+  return mapStableHloOpToStdScalarOp<stablehlo::DivOp>(
+      loc, resultTypes, resultTypes, {{one, oneAddExprNegX}}, b);
+}
+
+template <>
+inline Value mapStableHloOpToStdScalarOp<stablehlo::PowOp>(
+    Location loc, ArrayRef<Type> resultTypes, ArrayRef<Type> argTypes,
+    stablehlo::PowOp::Adaptor adaptor, OpBuilder* b) {
+  auto lb = ImplicitLocOpBuilder(loc, *b);
+  // Floating point can use std::powf
+  auto resultType = resultTypes.front();
+  if (resultType.isa<ComplexType, FloatType>()) {
+    return MapStableHloOpToScalarOpImpl<IsFloatType, math::PowFOp,
+                                        IsComplexType, complex::PowOp>{}(
+        loc, resultTypes, argTypes, adaptor.getOperands(), b);
+  }
+
+  // Exponentiation by squaring:
+  // https://en.wikipedia.org/wiki/Exponentiation_by_squaring;
+  Value negOne =
+      lb.create<arith::ConstantOp>(lb.getIntegerAttr(resultType, -1));
+  Value zero = lb.create<arith::ConstantOp>(lb.getIntegerAttr(resultType, 0));
+  Value one = lb.create<arith::ConstantOp>(lb.getIntegerAttr(resultType, 1));
+  Value two = lb.create<arith::ConstantOp>(lb.getIntegerAttr(resultType, 2));
+  Value step = lb.create<arith::ConstantIndexOp>(1);
+  Value lowerBound = lb.create<arith::ConstantIndexOp>(0);
+  // Everything else would overflow for any exponent > 1, as 2^64
+  // is the larget possible exponent for a 64-bit integer, and
+  // that's 1 << 6.
+  Value upperBound = lb.create<arith::ConstantIndexOp>(6);
+  auto originalBase = adaptor.getLhs();
+  auto originalExponent = adaptor.getRhs();
+
+  Value accum =
+      lb.create<scf::ForOp>(
+            lowerBound, upperBound, step,
+            SmallVector<Value>({one, originalBase, originalExponent}),
+            [&](OpBuilder& b, Location, Value /*v*/, ValueRange iters) {
+              Value accum = iters[0];
+              Value base = iters[1];
+              Value exponent = iters[2];
+
+              Value condition = b.create<arith::CmpIOp>(
+                  loc, arith::CmpIPredicate::eq,
+                  b.create<::mlir::arith::AndIOp>(loc, exponent, one), one);
+              Value multiplied =
+                  b.create<::mlir::arith::MulIOp>(loc, accum, base);
+              accum = b.create<::mlir::arith::SelectOp>(loc, condition,
+                                                        multiplied, accum);
+              base = b.create<::mlir::arith::MulIOp>(loc, base, base);
+              exponent = b.create<::mlir::arith::ShRUIOp>(loc, exponent, one);
+              b.create<scf::YieldOp>(
+                  loc, SmallVector<Value>({accum, base, exponent}));
+            })
+          .getResult(0);
+
+  Value rhsIsEven = lb.create<arith::CmpIOp>(
+      arith::CmpIPredicate::eq,
+      lb.create<arith::RemSIOp>(adaptor.getRhs(), two), zero);
+  Value rhsIsNegative = lb.create<arith::CmpIOp>(arith::CmpIPredicate::slt,
+                                                 adaptor.getRhs(), zero);
+  Value lhsIsOne =
+      lb.create<arith::CmpIOp>(arith::CmpIPredicate::eq, adaptor.getLhs(), one);
+  Value lhsIsNegOne = lb.create<arith::CmpIOp>(arith::CmpIPredicate::eq,
+                                               adaptor.getLhs(), negOne);
+
+  // The accum is correct when the rhs is non-negative. When rhs is
+  // negative, we return 0 for integer, with the exception of lhs values of 1
+  // and -1 which have integer results for negative exponents. Specifically, the
+  // calulation is the following:
+  //
+  // - Return accum if the rhs is not negative.
+  // - Return 1 or -1 depending on the parity of rhs when the lhs is -1.
+  // - Return 1 if lhs is 1.
+  // - Else return 0.
+  Value ifLhsIsOne = lb.create<::mlir::arith::SelectOp>(lhsIsOne, one, zero);
+  Value ifLhsIsNegOne = lb.create<::mlir::arith::SelectOp>(
+      lhsIsNegOne, lb.create<::mlir::arith::SelectOp>(rhsIsEven, one, negOne),
+      ifLhsIsOne);
+  return lb.create<::mlir::arith::SelectOp>(rhsIsNegative, ifLhsIsNegOne,
+                                            accum);
+}
+
+template <>
+inline Value mapStableHloOpToStdScalarOp<stablehlo::SelectOp>(
+    Location loc, ArrayRef<Type> resultTypes, ArrayRef<Type> argTypes,
+    stablehlo::SelectOp::Adaptor adaptor, OpBuilder* b) {
+  return MapStableHloOpToScalarOpImpl<::mlir::arith::SelectOp>{}(
+      loc, resultTypes, argTypes, adaptor.getOperands(), b);
+}
+
+template <>
+inline Value mapStableHloOpToStdScalarOp<stablehlo::SignOp>(
+    Location loc, ArrayRef<Type> resultTypes, ArrayRef<Type> /*argTypes*/,
+    stablehlo::SignOp::Adaptor adaptor, OpBuilder* b) {
+  Value operand = adaptor.getOperand();
+  Type elementType = getElementTypeOrSelf(operand.getType());
+  if (auto floatType = elementType.dyn_cast<FloatType>()) {
+    Value zero =
+        b->create<arith::ConstantOp>(loc, b->getZeroAttr(operand.getType()));
+    Value ne0I1 = b->create<::mlir::arith::CmpFOp>(
+        loc, arith::CmpFPredicate::ONE, operand, zero);
+    Value ne0Float =
+        b->create<::mlir::arith::UIToFPOp>(loc, zero.getType(), ne0I1);
+    Value copySign = b->create<::mlir::math::CopySignOp>(loc, resultTypes,
+                                                         ne0Float, operand);
+    auto isNan = b->create<::mlir::arith::CmpFOp>(
+        loc, arith::CmpFPredicate::UNO, operand, operand);
+    return b->create<::mlir::arith::SelectOp>(loc, isNan, operand, copySign);
+  }
+  if (auto integerType = elementType.dyn_cast<IntegerType>()) {
+    // sign(x) = x == 0 ? 0 : ((x s>> 31) | 1)
+    Value zero =
+        b->create<arith::ConstantOp>(loc, b->getZeroAttr(operand.getType()));
+    Value bitwidthMinusOne = getConstantOrSplat(
+        b, loc, operand.getType(),
+        b->getIntegerAttr(integerType, integerType.getWidth() - 1));
+    Value one = getConstantOrSplat(b, loc, operand.getType(),
+                                   b->getIntegerAttr(integerType, 1));
+    Value cmp = b->create<::mlir::arith::CmpIOp>(loc, arith::CmpIPredicate::eq,
+                                                 operand, zero);
+    Value ashr =
+        b->create<::mlir::arith::ShRSIOp>(loc, operand, bitwidthMinusOne);
+    Value orOp = b->create<::mlir::arith::OrIOp>(loc, ashr, one);
+    return b->create<::mlir::arith::SelectOp>(loc, cmp, zero, orOp);
+  }
+  if (elementType.isa<ComplexType>()) {
+    return b->create<::mlir::complex::SignOp>(loc, elementType, operand);
+  }
+  return nullptr;
+}
+
+/// Construct operations to select the saturated value if the shift amount is
+/// greater than the bitwidth of the type.
+inline Value selectShiftedOrSaturated(ImplicitLocOpBuilder& lb, Value rhs,
+                                      Value shifted, Value saturated,
+                                      Type type) {
+  Type etype =
+      type.isa<ShapedType>() ? type.cast<ShapedType>().getElementType() : type;
+  auto bitWidthInt = etype.getIntOrFloatBitWidth();
+  Value bitWidth = getConstantOrSplat(&lb, lb.getLoc(), type,
+                                      lb.getIntegerAttr(etype, bitWidthInt));
+  Value cmp = lb.create<mlir::arith::CmpIOp>(mlir::arith::CmpIPredicate::ugt,
+                                             bitWidth, rhs);
+  return lb.create<mlir::arith::SelectOp>(cmp, shifted, saturated);
+}
+
+template <>
+inline Value mapStableHloOpToStdScalarOp<stablehlo::ShiftLeftOp>(
+    Location loc, ArrayRef<Type> /*ResultTypes*/, ArrayRef<Type> /*argTypes*/,
+    stablehlo::ShiftLeftOp::Adaptor adaptor, OpBuilder* b) {
+  ImplicitLocOpBuilder lb(loc, *b);
+  Value lhs = adaptor.getLhs();
+  Value rhs = adaptor.getRhs();
+  Type type = lhs.getType();
+
+  // "Saturate" if the shift is greater than the bitwidth of the type
+  Value zero = lb.create<arith::ConstantOp>(lb.getZeroAttr(type));
+  Value shifted = lb.create<mlir::arith::ShLIOp>(lhs, rhs);
+
+  return selectShiftedOrSaturated(lb, rhs, shifted, zero, type);
+}
+
+template <>
+inline Value mapStableHloOpToStdScalarOp<stablehlo::ShiftRightLogicalOp>(
+    Location loc, ArrayRef<Type> /*ResultTypes*/, ArrayRef<Type> /*argTypes*/,
+    stablehlo::ShiftRightLogicalOp::Adaptor adaptor, OpBuilder* b) {
+  ImplicitLocOpBuilder lb(loc, *b);
+  Value lhs = adaptor.getLhs();
+  Value rhs = adaptor.getRhs();
+  Type type = lhs.getType();
+
+  // "Saturate" if the shift is greater than the bitwidth of the type
+  Value zero = lb.create<arith::ConstantOp>(b->getZeroAttr(type));
+  Value shifted = lb.create<mlir::arith::ShRUIOp>(lhs, rhs);
+
+  return selectShiftedOrSaturated(lb, rhs, shifted, zero, type);
+}
+
+template <>
+inline Value mapStableHloOpToStdScalarOp<stablehlo::ShiftRightArithmeticOp>(
+    Location loc, ArrayRef<Type> /*ResultTypes*/, ArrayRef<Type> /*argTypes*/,
+    stablehlo::ShiftRightArithmeticOp::Adaptor adaptor, OpBuilder* b) {
+  ImplicitLocOpBuilder lb(loc, *b);
+  Value lhs = adaptor.getLhs();
+  Value rhs = adaptor.getRhs();
+  Type type = lhs.getType();
+  Type etype =
+      type.isa<ShapedType>() ? type.cast<ShapedType>().getElementType() : type;
+  auto bitWidthInt = etype.getIntOrFloatBitWidth();
+
+  // "Saturate" if the shift is greater than the bitwidth of the type
+  Value maxShift = getConstantOrSplat(
+      b, loc, type, lb.getIntegerAttr(etype, bitWidthInt - 1));
+  Value saturatedShifted = lb.create<mlir::arith::ShRSIOp>(lhs, maxShift);
+  Value shifted = lb.create<mlir::arith::ShRSIOp>(lhs, rhs);
+
+  return selectShiftedOrSaturated(lb, rhs, shifted, saturatedShifted, type);
+}
+}  // namespace impl
+
+struct StableHloOpToStdScalarOp {
+  // Converts mhlo 'op' to linalg and arith ops.
+  template <typename StableHloOpTy>
+  static Value mapOp(StableHloOpTy op, ArrayRef<Type> resultTypes,
+                     ValueRange args, OpBuilder* b) {
+    auto argTypes = llvm::to_vector(op->getOperandTypes());
+    return mapOpWithArgTypes(op, resultTypes, argTypes, args, b);
+  }
+
+  // Converts mhlo 'op' to linalg and arith ops. The types of 'args' may already
+  // be converted, 'argTypes' are their original types.
+  template <typename StableHloOpTy>
+  static Value mapOpWithArgTypes(StableHloOpTy op, ArrayRef<Type> resultTypes,
+                                 ArrayRef<Type> argTypes, ValueRange args,
+                                 OpBuilder* b) {
+    static_assert(!std::is_same<StableHloOpTy, stablehlo::ConvertOp>::value);
+    return mapOpOfType<StableHloOpTy>(
+        op.getLoc(), resultTypes, argTypes,
+        typename StableHloOpTy::Adaptor(args, op->getAttrDictionary()), b);
+  }
+  // Overload for stablehlo::ConvertOp.
+  static Value mapOpWithArgTypes(stablehlo::ConvertOp op,
+                                 ArrayRef<Type> resultTypes,
+                                 ArrayRef<Type> argTypes, ValueRange args,
+                                 OpBuilder* b) {
+    return impl::mapConvertOpToStdScalarOp(op.getLoc(), op.getType(),
+                                           resultTypes, argTypes, args, b);
+  }
+
+  // Converts mhlo 'op' to linalg and arith ops.
+  template <typename StableHloOpTy>
+  static Value mapOpOfType(Location loc, ArrayRef<Type> resultTypes,
+                           ArrayRef<Type> argTypes,
+                           typename StableHloOpTy::Adaptor adaptor,
+                           OpBuilder* b) {
+    if (std::is_same<StableHloOpTy, stablehlo::ConvertOp>::value) {
+      // Note: this assumes that the caller is passing result/arg types with
+      // appropriate signedness.
+      return impl::mapConvertOpToStdScalarOp(
+          loc, resultTypes, resultTypes, argTypes, adaptor.getOperands(), b);
+    }
+    return impl::mapStableHloOpToStdScalarOp<StableHloOpTy>(
+        loc, resultTypes, argTypes, adaptor, b);
+  }
+};
+
+}  // namespace stablehlo
+}  // namespace mlir
+
+#endif  // IREE_COMPILER_INPUTCONVERSION_STABLEHLO_MAP_STABLEHLO_TO_SCALAR_OP_H
diff --git a/compiler/src/iree/compiler/InputConversion/StableHLO/PassDetail.h b/compiler/src/iree/compiler/InputConversion/StableHLO/PassDetail.h
new file mode 100644
index 000000000000..26e9cf31e8a2
--- /dev/null
+++ b/compiler/src/iree/compiler/InputConversion/StableHLO/PassDetail.h
@@ -0,0 +1,21 @@
+// Copyright 2019 The IREE Authors
+//
+// Licensed under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+#ifndef IREE_COMPILER_INPUTCONVERSION_STABLEHLO_PASSDETAIL_H_
+#define IREE_COMPILER_INPUTCONVERSION_STABLEHLO_PASSDETAIL_H_
+
+#include "mlir/Dialect/Func/IR/FuncOps.h"
+#include "mlir/IR/BuiltinOps.h"
+#include "mlir/Pass/Pass.h"
+
+namespace mlir::iree_compiler::stablehlo {
+
+#define GEN_PASS_DECL
+#include "iree/compiler/InputConversion/StableHLO/Passes.h.inc"
+
+}  // namespace mlir::iree_compiler::stablehlo
+
+#endif  // IREE_COMPILER_INPUTCONVERSION_STABLEHLO_PASSDETAIL_H_
diff --git a/compiler/src/iree/compiler/InputConversion/StableHLO/Passes.cpp b/compiler/src/iree/compiler/InputConversion/StableHLO/Passes.cpp
new file mode 100644
index 000000000000..6851d0b0cb9a
--- /dev/null
+++ b/compiler/src/iree/compiler/InputConversion/StableHLO/Passes.cpp
@@ -0,0 +1,25 @@
+// Copyright 2019 The IREE Authors
+//
+// Licensed under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+#include "iree/compiler/InputConversion/StableHLO/Passes.h"
+
+#include "mlir/Pass/PassManager.h"
+#include "mlir/Pass/PassOptions.h"
+#include "mlir/Pass/PassRegistry.h"
+
+namespace mlir::iree_compiler::stablehlo {
+
+namespace {
+#define GEN_PASS_REGISTRATION
+#include "iree/compiler/InputConversion/StableHLO/Passes.h.inc"  // IWYU pragma: export
+}  // namespace
+
+void registerStableHLOConversionPasses() {
+  // Generated.
+  registerPasses();
+}
+
+}  // namespace mlir::iree_compiler::stablehlo
diff --git a/compiler/src/iree/compiler/InputConversion/StableHLO/Passes.h b/compiler/src/iree/compiler/InputConversion/StableHLO/Passes.h
new file mode 100644
index 000000000000..d3420701e23d
--- /dev/null
+++ b/compiler/src/iree/compiler/InputConversion/StableHLO/Passes.h
@@ -0,0 +1,33 @@
+// Copyright 2019 The IREE Authors
+//
+// Licensed under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+#ifndef IREE_COMPILER_INPUTCONVERSION_STABLEHLO_PASSES_H_
+#define IREE_COMPILER_INPUTCONVERSION_STABLEHLO_PASSES_H_
+
+#include "iree/compiler/InputConversion/StableHLO/PassDetail.h"
+#include "mlir/Pass/Pass.h"
+
+namespace mlir {
+class TypeConverter;
+namespace iree_compiler::stablehlo {
+
+std::unique_ptr<TypeConverter> createStableHloToLinalgTypeConverter();
+
+//===----------------------------------------------------------------------===//
+// Pipelines
+//===----------------------------------------------------------------------===//
+
+// TODO: Add a full conversion pipeline and passes.
+
+//===----------------------------------------------------------------------===//
+// Register all Passes
+//===----------------------------------------------------------------------===//
+
+void registerStableHLOConversionPasses();
+
+}  // namespace iree_compiler::stablehlo
+}  // namespace mlir
+#endif  // IREE_COMPILER_INPUTCONVERSION_STABLEHLO_PASSES_H_
diff --git a/compiler/src/iree/compiler/InputConversion/StableHLO/Passes.td b/compiler/src/iree/compiler/InputConversion/StableHLO/Passes.td
new file mode 100644
index 000000000000..e2a7256ed177
--- /dev/null
+++ b/compiler/src/iree/compiler/InputConversion/StableHLO/Passes.td
@@ -0,0 +1,21 @@
+// Copyright 2019 The IREE Authors
+//
+// Licensed under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+#ifndef IREE_COMPILER_INPUTCONVERSION_STABLEHLO_PASSES
+#define IREE_COMPILER_INPUTCONVERSION_STABLEHLO_PASSES
+
+include "mlir/Pass/PassBase.td"
+
+def ConvertStableHloToLinalg :
+    Pass<"iree-stablehlo-to-linalg", "ModuleOp"> {
+  let summary = "Convert from StableHLO ops to Linalg ops on tensors";
+  let options = [Option<"enablePrimitiveOps", "enable-primitive-ops", "bool",
+                        /*default=*/"false",
+                        "Lower to primitive Linalg ops (map, reduce and "
+                        "transpose) when possible, instead of linalg.generic">];
+}
+
+#endif // IREE_COMPILER_INPUTCONVERSION_STABLEHLO_PASSES
diff --git a/compiler/src/iree/compiler/InputConversion/StableHLO/Rewriters.h b/compiler/src/iree/compiler/InputConversion/StableHLO/Rewriters.h
new file mode 100644
index 000000000000..36cea233fb3f
--- /dev/null
+++ b/compiler/src/iree/compiler/InputConversion/StableHLO/Rewriters.h
@@ -0,0 +1,22 @@
+// Copyright 2019 The IREE Authors
+//
+// Licensed under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+#ifndef IREE_COMPILER_INPUTCONVERSION_STABLEHLO_REWRITERS_H_
+#define IREE_COMPILER_INPUTCONVERSION_STABLEHLO_REWRITERS_H_
+
+#include "mlir/Transforms/DialectConversion.h"
+
+namespace mlir::iree_compiler::stablehlo {
+
+/// Populates the patterns that convert from StableHLO to Linalg on tensors.
+void populateStableHloToLinalgConversionPatterns(MLIRContext* context,
+                                                 TypeConverter& typeConverter,
+                                                 RewritePatternSet* patterns,
+                                                 bool enablePrimitiveOps);
+
+}  // namespace mlir::iree_compiler::stablehlo
+
+#endif  // IREE_COMPILER_INPUTCONVERSION_STABLEHLO_REWRITERS_H_
diff --git a/compiler/src/iree/compiler/InputConversion/StableHLO/StableHLOToLinalg.cpp b/compiler/src/iree/compiler/InputConversion/StableHLO/StableHLOToLinalg.cpp
new file mode 100644
index 000000000000..2e8cec3ea0c0
--- /dev/null
+++ b/compiler/src/iree/compiler/InputConversion/StableHLO/StableHLOToLinalg.cpp
@@ -0,0 +1,3246 @@
+// Copyright 2019 The IREE Authors
+//
+// Licensed under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+// Implements logic for lowering StableHLO/CHLO dialects to Linalg dialect.
+
+#include <algorithm>
+#include <cstdint>
+#include <memory>
+#include <numeric>
+#include <optional>
+#include <string>
+#include <utility>
+
+#include "iree/compiler/InputConversion/StableHLO/LegalizeToLinalgUtils.h"
+#include "iree/compiler/InputConversion/StableHLO/Passes.h"
+#include "iree/compiler/InputConversion/StableHLO/Rewriters.h"
+#include "iree/compiler/InputConversion/StableHLO/TypeConversion.h"
+#include "llvm/ADT/ArrayRef.h"
+#include "llvm/ADT/SmallVector.h"
+#include "mlir/Dialect/Arith/IR/Arith.h"
+#include "mlir/Dialect/Bufferization/IR/Bufferization.h"
+#include "mlir/Dialect/Complex/IR/Complex.h"
+#include "mlir/Dialect/Func/IR/FuncOps.h"
+#include "mlir/Dialect/Linalg/IR/Linalg.h"
+#include "mlir/Dialect/Linalg/Transforms/Transforms.h"
+#include "mlir/Dialect/Linalg/Utils/Utils.h"
+#include "mlir/Dialect/Math/IR/Math.h"
+#include "mlir/Dialect/MemRef/IR/MemRef.h"
+#include "mlir/Dialect/SCF/IR/SCF.h"
+#include "mlir/Dialect/Shape/IR/Shape.h"
+#include "mlir/Dialect/SparseTensor/IR/SparseTensor.h"
+#include "mlir/Dialect/Tensor/IR/Tensor.h"
+#include "mlir/Dialect/Utils/StructuredOpsUtils.h"
+#include "mlir/IR/AffineExpr.h"
+#include "mlir/IR/Attributes.h"
+#include "mlir/IR/Builders.h"
+#include "mlir/IR/BuiltinAttributes.h"
+#include "mlir/IR/BuiltinOps.h"
+#include "mlir/IR/BuiltinTypes.h"
+#include "mlir/IR/Location.h"
+#include "mlir/IR/MLIRContext.h"
+#include "mlir/IR/Operation.h"
+#include "mlir/IR/PatternMatch.h"
+#include "mlir/IR/TypeUtilities.h"
+#include "mlir/Pass/Pass.h"
+#include "mlir/Support/LLVM.h"
+#include "mlir/Support/LogicalResult.h"
+#include "mlir/Transforms/DialectConversion.h"
+#include "stablehlo/dialect/StablehloOps.h"
+
+namespace mlir::iree_compiler::stablehlo {
+
+#define GEN_PASS_DEF_CONVERTSTABLEHLOTOLINALG
+#include "iree/compiler/InputConversion/StableHLO/Passes.h.inc"
+
+namespace {
+
+namespace stablehlo = mlir::stablehlo;
+
+Value getResultValue(Operation* op) { return op->getResult(0); }
+
+ShapedType getHloOpResultType(Operation* op) {
+  return getResultValue(op).getType().cast<ShapedType>();
+}
+
+bool verifyHloOpBufferOrTensorSemantics(Operation* op) {
+  auto verifyType = [&](Value val) -> bool {
+    return val.getType().isa<RankedTensorType>();
+  };
+  if (!llvm::all_of(op->getOperands(), verifyType)) return false;
+  return llvm::all_of(op->getResults(), verifyType);
+}
+
+Value fillTensorWithZeros(OpBuilder& builder, Location loc, Value tensor) {
+  auto type = tensor.getType().cast<ShapedType>();
+  Value zero;
+  // Complex numbers are a special case.
+  if (auto complexType = type.getElementType().dyn_cast<ComplexType>()) {
+    auto zeroElement = builder.getZeroAttr(complexType.getElementType());
+    auto zeroAttr = builder.getArrayAttr({zeroElement, zeroElement});
+    zero = builder.create<complex::ConstantOp>(loc, complexType, zeroAttr);
+  } else {
+    auto zeroAttr = builder.getZeroAttr(type.getElementType());
+    zero = builder.create<arith::ConstantOp>(loc, zeroAttr);
+  }
+  return builder.create<linalg::FillOp>(loc, zero, tensor).result();
+}
+
+SmallVector<int64_t, 4> extract1DVector(DenseIntElementsAttr elements) {
+  SmallVector<int64_t, 4> ret;
+  for (const APInt& element : elements) {
+    ret.push_back(element.getLimitedValue());
+  }
+  return ret;
+}
+
+/// Returns a permutation AffineMap that puts all reduction dimensions to the
+/// last. The order of parallel loops and reduction loops are all sorted. E.g.,
+/// if `rank` is 4 and `reductionDims` is {1, 3}, then
+/// "(d0, d1, d2, d3) -> (d0, d2, d1, d3)" is used. The inverse permutation of
+/// the AffineMap is returned.
+AffineMap getTransposeMapForReduction(MLIRContext* context, int rank,
+                                      ArrayRef<int64_t> reductionDims) {
+  llvm::SmallSetVector<int, 4> s;
+  for (auto dim : reductionDims) s.insert(dim);
+
+  SmallVector<unsigned, 4> permutation;
+  for (int i = 0; i < rank; ++i)
+    if (!s.count(i)) permutation.push_back(i);
+  for (auto dim : reductionDims) permutation.push_back(dim);
+
+  auto map = AffineMap::getPermutationMap(permutation, context);
+  return inversePermutation(map);
+}
+
+/// Returns true if the given `attr` is a splat of the given `value`.
+bool isSplatValue(DenseIntElementsAttr attr, uint64_t value) {
+  return attr.isSplat() && attr.getSplatValue<uint64_t>() == value;
+}
+
+/// Extracts an element from a tensor and optionally converts it to an index
+/// type, based on the tensor's pre-type conversion type.
+Value extractIndexFromTensor(OpBuilder& builder, Location loc, Value tensor,
+                             ShapedType originalType,
+                             ArrayRef<Value> tensorIndex = {}) {
+  Value extracted = builder.create<tensor::ExtractOp>(loc, tensor, tensorIndex);
+  if (extracted.getType().isIndex()) return extracted;
+  return originalType.getElementType().isUnsignedInteger()
+             ? builder.createOrFold<arith::IndexCastUIOp>(
+                   loc, builder.getIndexType(), extracted)
+             : builder.createOrFold<arith::IndexCastOp>(
+                   loc, builder.getIndexType(), extracted);
+}
+
+/// Ensures a tensor has the same shape (not including the element type) as
+/// another.
+Value coerceTensorShape(OpBuilder& builder, Location loc,
+                        TypedValue<ShapedType> value, ShapedType targetType) {
+  return builder.createOrFold<tensor::CastOp>(
+      loc, targetType.cloneWith(std::nullopt, value.getType().getElementType()),
+      value);
+}
+
+//===----------------------------------------------------------------------===//
+// stablehlo.RngOp conversion patterns.
+//===----------------------------------------------------------------------===//
+
+// Pass to lower from rng to stateless pseudo RNG with LCG
+// algorithm
+struct RngUniformConversion : public OpConversionPattern<stablehlo::RngOp> {
+  using OpConversionPattern<stablehlo::RngOp>::OpConversionPattern;
+
+  LogicalResult matchAndRewrite(
+      stablehlo::RngOp op, OpAdaptor adaptor,
+      ConversionPatternRewriter& rewriter) const final {
+    // We only handle uniform distributions
+    if (op.getRngDistribution() !=
+        ::mlir::stablehlo::RngDistribution::UNIFORM) {
+      return failure();
+    }
+    // TODO(raikonenfnu): Handle other element types as well.
+    auto minTy = adaptor.getOperands()[0].getType().dyn_cast<ShapedType>();
+    auto maxTy = adaptor.getOperands()[0].getType().dyn_cast<ShapedType>();
+    if (!minTy.getElementType().dyn_cast<FloatType>() ||
+        !maxTy.getElementType().dyn_cast<FloatType>()) {
+      return rewriter.notifyMatchFailure(
+          op, "expected min/max for rng op to be FloatType");
+    }
+    auto targetTy = this->typeConverter->convertType(op.getResult().getType())
+                        .cast<ShapedType>();
+    if (!targetTy) {
+      return rewriter.notifyMatchFailure(
+          op, "expected target shape of rng op to be ShapedType");
+    }
+    auto loc = op.getLoc();
+    Value emptyTensor =
+        getEmptyTensorFor(rewriter, loc, targetTy, op, adaptor.getOperands());
+    // Creates index map using target matrix's rank.
+    auto targetRank = targetTy.getRank();
+    SmallVector<AffineMap, 3> indexingMaps(
+        2, AffineMap::get(targetRank, /*symbolCount=*/0,
+                          SmallVector<AffineExpr>({}), rewriter.getContext()));
+    indexingMaps.push_back(rewriter.getMultiDimIdentityMap(targetRank));
+    const int kInitialSeed = 0;
+    // Generic region with LCG Algorithm that make use of element index from:
+    // https://reviews.llvm.org/D101364
+    auto linalgOp = rewriter.create<linalg::GenericOp>(
+        loc, /*resultTensors=*/targetTy,
+        /*inputs=*/
+        ValueRange{adaptor.getOperands()[0], adaptor.getOperands()[1]},
+        /*outputs=*/emptyTensor, indexingMaps,
+        getParallelAndReductionIterators(/*nLoops=*/targetRank,
+                                         /*nReduction=*/0),
+        [&](OpBuilder& b, Location loc, ValueRange args) {
+          llvm::SmallVector<Value> updateVec = {b.create<arith::ConstantOp>(
+              loc, b.getI32IntegerAttr(kInitialSeed))};
+          Value multiplier =
+              b.create<arith::ConstantOp>(loc, b.getI32IntegerAttr(1103515245));
+          Value incrementStep =
+              b.create<arith::ConstantOp>(loc, b.getI32IntegerAttr(12345));
+          // For output matrix with rank N:
+          // temp1 = (cast(I32, index(D.0)) + seed) * mult + incr
+          // ...
+          // tempN = (cast(I32, index(D.(N))) + tempN_1) * mult + incr
+          for (int i = 0; i < targetRank; i++) {
+            Value update = updateVec.back();
+            Value ind = b.create<linalg::IndexOp>(loc, i);
+            Value castInd =
+                b.create<arith::IndexCastOp>(loc, b.getI32Type(), ind);
+            Value addRes = b.create<arith::AddIOp>(loc, castInd, update);
+            Value multRes = b.create<arith::MulIOp>(loc, addRes, multiplier);
+            Value incRes = b.create<arith::AddIOp>(loc, multRes, incrementStep);
+            updateVec.push_back(incRes);
+          }
+          // Scaling = (max - min) * const(F64, 2.3283064E-10)
+          // which is derived from rand(min,max) = rand()/(RAND_MAX/(max-min)).
+          Value epsilon = b.create<arith::ConstantOp>(
+              loc, b.getFloatAttr(args[0].getType(), 2.3283064E-10));
+          Value range = b.create<arith::SubFOp>(loc, args[1], args[0]);
+          Value scale = b.create<arith::MulFOp>(loc, range, epsilon);
+          // Res = cast(T, cast(F64, tempN) * scaling + min)
+          Value updateCast = b.create<arith::UIToFPOp>(
+              loc, targetTy.getElementType(), updateVec.back());
+          Value scaleUpdate = b.create<arith::MulFOp>(loc, updateCast, scale);
+          Value res = b.create<arith::AddFOp>(loc, scaleUpdate, args[0]);
+          b.create<linalg::YieldOp>(loc, res);
+        },
+        linalg::getPrunedAttributeList(op));
+    rewriter.replaceOp(op, linalgOp.getResults());
+    return success();
+  }
+};
+
+//===----------------------------------------------------------------------===//
+// stablehlo.Einsum conversion patterns.
+//===----------------------------------------------------------------------===//
+
+// Looks through a set of dimension that has been marked as reduction axes,
+// if it is found within the set, then we set it as "reduction", otherwise
+// we can label it as "parallel".
+SmallVector<utils::IteratorType, 3> getEinsumLoopsAttrs(
+    const llvm::SmallSetVector<StringRef, 4>& inputInd,
+    const llvm::SmallSetVector<StringRef, 4>& reductionDims) {
+  SmallVector<utils::IteratorType, 3> res;
+  for (StringRef dim : inputInd) {
+    if (!reductionDims.contains(dim)) {
+      res.push_back(utils::IteratorType::parallel);
+    } else {
+      res.push_back(utils::IteratorType::reduction);
+    }
+  }
+  return res;
+}
+
+SmallVector<Value, 2> extractDynamicEinsumSizes(
+    OpBuilder& b, Location loc, Value lhs, Value rhs,
+    const SmallVector<std::string>& lhsLoopVec,
+    const SmallVector<std::string>& rhsLoopVec,
+    const SmallVector<std::string>& outputLoopVec) {
+  SmallVector<Value, 2> dynSizes;
+  for (const std::string& dimInd : outputLoopVec) {
+    Value dimSize;
+    const auto* dimIndIt =
+        std::find(lhsLoopVec.begin(), lhsLoopVec.end(), dimInd);
+    if (dimIndIt != lhsLoopVec.end()) {
+      // Query from lhs vars.
+      auto dimIndPos = dimIndIt - lhsLoopVec.begin();
+      auto lhsShape = lhs.getType().dyn_cast<RankedTensorType>().getShape();
+      if (lhsShape[dimIndPos] != ShapedType::kDynamic) continue;
+      dimSize = b.create<tensor::DimOp>(loc, lhs, dimIndPos);
+    } else {
+      // query from rhs vars.
+      dimIndIt = std::find(rhsLoopVec.begin(), rhsLoopVec.end(), dimInd);
+      auto dimIndPos = dimIndIt - rhsLoopVec.begin();
+      auto rhsShape = rhs.getType().dyn_cast<RankedTensorType>().getShape();
+      if (rhsShape[dimIndPos] != ShapedType::kDynamic) continue;
+      dimSize = b.create<tensor::DimOp>(loc, rhs, dimIndPos);
+    }
+    dynSizes.push_back(dimSize);
+  }
+  return dynSizes;
+}
+
+// Adds indices/axes that are missing from output set.
+llvm::SmallSetVector<StringRef, 4> findSummationAxes(
+    const llvm::SmallSetVector<StringRef, 4>& inputSet,
+    const llvm::SmallSetVector<StringRef, 4>& outputSet) {
+  llvm::SmallSetVector<StringRef, 4> summationAxes;
+  for (StringRef ind : inputSet) {
+    if (!outputSet.contains(ind)) summationAxes.insert(ind);
+  }
+  return summationAxes;
+}
+
+// Given a 1:1 map from std::string -> affine dimension expression
+// we can get the affine expression of dimensions that an
+// operand will access based on the input_str of einsum_config.
+// For example:
+// let string_dim_umap = {'a' : d0, 'b' : d1, 'c' : d2}
+// for einsum_config "abc,cb->acb"
+// first_input_operand will get umap[{"a","b","c"}] -> (d0, d1, d2).
+// second_input_operand will get umap[{"c","b"}] -> (d2, d1).
+// output_operand will get umap[{"a","c","b"}] -> (d0, d2, d1).
+SmallVector<AffineExpr> getExprFromConfig(
+    const SmallVector<std::string>& loopDims,
+    const DenseMap<StringRef, AffineExpr>& strAffineDimUmap) {
+  SmallVector<AffineExpr> exprs;
+  for (const auto& dim : loopDims) {
+    exprs.push_back(strAffineDimUmap.lookup(dim));
+  }
+  return exprs;
+}
+
+// Convert mhlo.einsum op into linalg.generic.
+// Algorithm in general 3 steps:
+
+// Step1) Dissect entire einsum_config to different operands
+// e.g f("abc,cd->abd") = {lhs:["abc"], rhs:["cd"], out:["abd"]}.
+
+// Step2) Split up the string into vector of the elements
+// e.g {lhs:["abc"], rhs:["cd"], out:["abd"]} = {lhs:["a","b","c"],
+// rhs:["c","d"], out:["a","b","d"]}.
+
+// Step3) Convert the vector into data access
+// patern represented by affineMaps with affineDimensions e.g
+// {lhs:["a","b","c"], rhs:["c","d"], out:["a","b","d"]} = {lhs:[d0,d1,d2],
+// rhs:[d2,d3], out:[d0,d1,d3]}.
+class EinsumToLinalgConverter
+    : public OpConversionPattern<stablehlo::EinsumOp> {
+ public:
+  using OpConversionPattern<stablehlo::EinsumOp>::OpConversionPattern;
+
+  LogicalResult matchAndRewrite(
+      stablehlo::EinsumOp op, OpAdaptor adaptor,
+      ConversionPatternRewriter& rewriter) const final {
+    auto getRank = [](Value v) {
+      return v.getType().cast<ShapedType>().getRank();
+    };
+    auto einsumConfig = op.getEinsumConfig();
+
+    // With the assumption of binary input operand and single output
+    // get the inputs and output operands' indices.
+    // einsum_config = "lhs_loop,rhs_loop->out_loop"
+    std::size_t posArrow = einsumConfig.find(kArrow);
+    std::size_t posComma = einsumConfig.find(kComma);
+
+    StringRef lhsLoop = einsumConfig.substr(0, posComma);
+    StringRef rhsLoop = einsumConfig.substr(
+        posComma + kComma.size(), posArrow - (posComma + kComma.size()));
+    StringRef outLoop = einsumConfig.substr(posArrow + kArrow.size());
+
+    // Check for Invalid Configs.
+    // 1.Check that there is only maximum 2 inputs
+    // 2.Check that there is only maximum 1 output
+    // 3.Check that there is 1 kArrow
+    if (rhsLoop.find(kComma) != std::string::npos ||
+        outLoop.find(kComma) != std::string::npos ||
+        outLoop.find(kArrow) != std::string::npos) {
+      return rewriter.notifyMatchFailure(op, "Invalid einsum config!");
+    }
+
+    // Find result type, if on tensors.
+    auto resultTy = this->typeConverter->convertType(getHloOpResultType(op))
+                        .dyn_cast<RankedTensorType>();
+
+    // Check result type compatibility.
+    if (!resultTy || !(resultTy.getElementType().isSignlessIntOrFloat())) {
+      return rewriter.notifyMatchFailure(op, "Invalid result type");
+    }
+
+    // Convert the representation to vector<string>.
+    SmallVector<std::string> lhsEin =
+        getEinsumConfigAsVector(lhsLoop, getRank(adaptor.getLhs()));
+    SmallVector<std::string> rhsEin =
+        getEinsumConfigAsVector(rhsLoop, getRank(adaptor.getRhs()));
+    SmallVector<std::string> outEin =
+        getEinsumConfigAsVector(outLoop, resultTy.getRank());
+
+    if (!checkBatchHasEqualRank(lhsEin.size(), lhsLoop, rhsEin.size(), rhsLoop,
+                                outEin.size(), outLoop)) {
+      return rewriter.notifyMatchFailure(
+          op, "Invalid elipsis('...') within einsum config!");
+    }
+
+    // Find all unique indices in the input and output.
+    llvm::SmallSetVector<StringRef, 4> inputInd;
+    llvm::SmallSetVector<StringRef, 4> outputInd;
+
+    inputInd.insert(lhsEin.begin(), lhsEin.end());
+    inputInd.insert(rhsEin.begin(), rhsEin.end());
+    outputInd.insert(outEin.begin(), outEin.end());
+
+    llvm::SmallSetVector<StringRef, 4> reductionAxe =
+        findSummationAxes(inputInd, outputInd);
+
+    // Find input/output values and types.
+    auto loc = op.getLoc();
+
+    // Prepare init tensor for linalg.generic op.
+    auto dynSizes =
+        extractDynamicEinsumSizes(rewriter, loc, adaptor.getLhs(),
+                                  adaptor.getRhs(), lhsEin, rhsEin, outEin);
+    Value output = getEmptyTensor(rewriter, loc, resultTy, dynSizes);
+    if (!reductionAxe.empty()) {
+      output = fillTensorWithZeros(rewriter, loc, output);
+    }
+
+    // Create indexing maps.
+    // Create a 1:1 map from f:strDimension -> affineDimension.
+    int64_t nloops = inputInd.size();
+    DenseMap<StringRef, AffineExpr> strAffineDimUmap;
+    for (const auto& it : llvm::enumerate(inputInd)) {
+      strAffineDimUmap[it.value()] = rewriter.getAffineDimExpr(it.index());
+    }
+
+    // From einsum_config of each operand in vector<string>, generate
+    // the equivalent vector<AffineExpr>.
+    SmallVector<AffineMap, 4> maps;
+    for (const SmallVector<std::string>& loopOperand :
+         {lhsEin, rhsEin, outEin}) {
+      auto exprs = getExprFromConfig(loopOperand, strAffineDimUmap);
+      maps.push_back(AffineMap::get(nloops, 0, exprs, rewriter.getContext()));
+    }
+
+    auto linalgOp = rewriter.create<linalg::GenericOp>(
+        loc, resultTy ? resultTy : TypeRange{}, adaptor.getOperands(), output,
+        maps, getEinsumLoopsAttrs(inputInd, reductionAxe),
+        [reductionAxe](OpBuilder& b, Location nestedLoc, ValueRange args) {
+          Value resultVal =
+              b.create<mlir::arith::MulFOp>(nestedLoc, args[0], args[1]);
+          if (!reductionAxe.empty()) {
+            resultVal =
+                b.create<mlir::arith::AddFOp>(nestedLoc, args[2], resultVal);
+          }
+          b.create<linalg::YieldOp>(nestedLoc, resultVal);
+        },
+        linalg::getPrunedAttributeList(op));
+    rewriter.replaceOp(op, linalgOp.getResults());
+    return success();
+  }
+
+ private:
+  static constexpr StringRef kArrow = "->";
+  static constexpr StringRef kComma = ",";
+  static constexpr StringRef kEllipsis = "...";
+
+  static bool checkBatchHasEqualRank(size_t lhsRank, StringRef lhsLoop,
+                                     size_t rhsRank, StringRef rhsLoop,
+                                     size_t outRank, StringRef outLoop);
+  static SmallVector<std::string> getEinsumConfigAsVector(StringRef loop,
+                                                          size_t operandRank);
+};
+
+// Definition of util const member variables.
+constexpr StringRef EinsumToLinalgConverter::kArrow;
+constexpr StringRef EinsumToLinalgConverter::kComma;
+constexpr StringRef EinsumToLinalgConverter::kEllipsis;
+
+// Convert the representation from string/vector<char> to vector<string>.
+// i.e ("abc") -> {"a", "b", "c"}. For cases with ellipsis with batch rank 3:
+// get loop_dim = f("ab...cde") = {"a","b","0","1","2","c","d","e"}
+SmallVector<std::string> EinsumToLinalgConverter::getEinsumConfigAsVector(
+    StringRef loop, size_t operandRank) {
+  SmallVector<std::string> loopDim;
+  size_t preElip = loop.find(kEllipsis);
+  bool hasElip = preElip != std::string::npos;
+  if (!hasElip) preElip = loop.size();
+  // Add the dimension until the end or up to ellipsis if it exist.
+  for (int64_t preElipInd = 0; preElipInd < static_cast<int64_t>(preElip);
+       preElipInd++) {
+    loopDim.push_back(loop.substr(preElipInd, 1).str());
+  }
+  if (!hasElip) return loopDim;
+  // Case where Ellipsis presence:
+  size_t nonBatchRank = loop.size() - kEllipsis.size();
+  size_t batchRank = operandRank - nonBatchRank;
+  // Add the batch dimension ("0",...,"N") where N is rank of batch into the
+  // loop.
+  for (int64_t batchInd = 0; batchInd < static_cast<int64_t>(batchRank);
+       batchInd++) {
+    loopDim.push_back(std::to_string(batchInd));
+  }
+  // Add the dimension after ellipsis into the loop.
+  int postElip = preElip + kEllipsis.size();
+  for (int64_t postElipInd = postElip;
+       postElipInd < static_cast<int64_t>(loop.size()); postElipInd++) {
+    loopDim.push_back(loop.substr(postElipInd, 1).str());
+  }
+  return loopDim;
+}
+
+// Returns true if all operand's batch has same rank.
+bool EinsumToLinalgConverter::checkBatchHasEqualRank(
+    size_t lhsRank, StringRef lhsLoop, size_t rhsRank, StringRef rhsLoop,
+    size_t outRank, StringRef outLoop) {
+  SmallVector<int, 3> batchRankVec;
+  if (lhsRank != lhsLoop.size()) {
+    size_t lhsBatchRank = lhsRank - (lhsLoop.size() - kEllipsis.size());
+    batchRankVec.push_back(lhsBatchRank);
+  }
+  if (rhsRank != rhsLoop.size()) {
+    size_t rhsBatchRank = rhsRank - (rhsLoop.size() - kEllipsis.size());
+    batchRankVec.push_back(rhsBatchRank);
+  }
+  if (outRank != outLoop.size()) {
+    size_t outBatchRank = outRank - (outLoop.size() - kEllipsis.size());
+    batchRankVec.push_back(outBatchRank);
+  }
+  bool batchHasEqualRank = true;
+
+  // Condition is valid if only 1 operand or less have batches.
+  if (batchRankVec.size() < 2) return batchHasEqualRank;
+  if (!std::equal(batchRankVec.begin() + 1, batchRankVec.end(),
+                  batchRankVec.begin()) &&
+      batchRankVec.size() > 1)
+    batchHasEqualRank = false;
+  return batchHasEqualRank;
+}
+
+/// Base class for lowering HLO operations that have one operand and one result,
+/// and are semantically equivalent to a copy of the input to the output (like
+/// transpose, some reshape, etc.). The derived classes need to provide a method
+/// `getIndexingMaps` that returns AffineMaps for the index maps of the input
+/// and the output.
+template <typename Derived, typename OpTy>
+class DataMovementOpConverter : public OpConversionPattern<OpTy> {
+ public:
+  using OpConversionPattern<OpTy>::OpConversionPattern;
+
+  LogicalResult matchAndRewrite(
+      OpTy op, typename OpTy::Adaptor adaptor,
+      ConversionPatternRewriter& rewriter) const final {
+    if (!verifyHloOpBufferOrTensorSemantics(op)) return failure();
+    auto resultType = getHloOpResultType(op);
+    resultType = this->typeConverter->convertType(resultType)
+                     .template cast<ShapedType>();
+
+    SmallVector<AffineMap, 2> indexingMaps =
+        Derived::getIndexingMaps(op, &rewriter);
+    if (indexingMaps.empty()) return failure();
+
+    auto nloops = resultType.getRank();
+    auto loc = op.getLoc();
+    auto linalgOp = rewriter.create<linalg::GenericOp>(
+        loc,
+        /*resultTensorTypes=*/resultType,
+        /*inputs=*/adaptor.getOperands().front(),
+        /*outputBuffers=*/
+
+        ValueRange{getEmptyTensorFor(rewriter, loc, resultType, op,
+                                     adaptor.getOperands())},
+        indexingMaps, getNParallelLoopsAttrs(nloops),
+        [&](OpBuilder& nestedBuilder, Location /*nested_loc*/,
+            ValueRange args) {
+          nestedBuilder.create<linalg::YieldOp>(loc, *args.begin());
+        },
+        linalg::getPrunedAttributeList(op));
+    rewriter.replaceOp(op, linalgOp.getOperation()->getResults());
+    return success();
+  }
+};
+
+/// Pattern to convert BroadcastOp to Linalg ops.
+template <typename OpTy>
+class BroadcastConverter
+    : public DataMovementOpConverter<BroadcastConverter<OpTy>, OpTy> {
+ public:
+  using DataMovementOpConverter<BroadcastConverter,
+                                OpTy>::DataMovementOpConverter;
+
+  static SmallVector<AffineMap, 2> getIndexingMaps(OpTy broadcastOp,
+                                                   Builder* b) {
+    ShapedType inputType =
+        broadcastOp.getOperand().getType().template cast<ShapedType>();
+    unsigned inputRank = inputType.getRank();
+    unsigned nloops = getHloOpResultType(broadcastOp).getRank();
+
+    // BroadcastOp prepends the dimensions in the `broadcast_sizes` attribute to
+    // the input's dimensions.
+    unsigned numPrependedDims = llvm::size(broadcastOp.getBroadcastSizes());
+    SmallVector<AffineExpr, 4> inputDimExprs;
+    inputDimExprs.reserve(inputRank);
+    for (unsigned i = 0; i < inputRank; ++i) {
+      inputDimExprs.push_back(b->getAffineDimExpr(numPrependedDims + i));
+    }
+
+    AffineMap inputMap;
+    MLIRContext* context = b->getContext();
+    if (inputDimExprs.empty()) {
+      // The input is a scalar, i.e. this is a scalar broadcast op.
+      inputMap = AffineMap::get(nloops, /*symbolCount=*/0, context);
+    } else {
+      inputMap =
+          AffineMap::get(nloops, /*symbolCount=*/0, inputDimExprs, context);
+    }
+    return {inputMap, b->getMultiDimIdentityMap(nloops)};
+  }
+};
+
+class BroadcastOpToBroadcastConverter
+    : public OpConversionPattern<stablehlo::BroadcastOp> {
+  using OpConversionPattern<stablehlo::BroadcastOp>::OpConversionPattern;
+
+  LogicalResult matchAndRewrite(
+      stablehlo::BroadcastOp op, OpAdaptor adaptor,
+      ConversionPatternRewriter& rewriter) const override {
+    auto resultTy = typeConverter->convertType(op.getType()).cast<ShapedType>();
+
+    int64_t numPrependedDims = op.getBroadcastSizes().size();
+    SmallVector<int64_t> dimensions =
+        llvm::to_vector(llvm::seq<int64_t>(0, numPrependedDims));
+
+    auto loc = op.getLoc();
+    Value emptyTensor =
+        getEmptyTensorFor(rewriter, loc, resultTy, op, adaptor.getOperands());
+
+    rewriter.replaceOpWithNewOp<linalg::BroadcastOp>(
+        op, op.getOperand(), emptyTensor, dimensions,
+        linalg::getPrunedAttributeList(op));
+    return success();
+  }
+};
+
+class HloBroadcastInDimConverter
+    : public DataMovementOpConverter<HloBroadcastInDimConverter,
+                                     stablehlo::BroadcastInDimOp> {
+ public:
+  using DataMovementOpConverter<
+      HloBroadcastInDimConverter,
+      stablehlo::BroadcastInDimOp>::DataMovementOpConverter;
+
+  static SmallVector<AffineMap, 2> getIndexingMaps(
+      stablehlo::BroadcastInDimOp broadcastOp, Builder* b) {
+    auto resultType = getHloOpResultType(broadcastOp);
+    auto operandType =
+        broadcastOp.getOperand().getType().template cast<ShapedType>();
+    unsigned nloops = resultType.getRank();
+
+    // The input is a scalar, i.e. this is a scalar broadcast op.
+    if (operandType.getRank() == 0) {
+      return {AffineMap::get(nloops, /*symbolCount=*/0, b->getContext()),
+              b->getMultiDimIdentityMap(nloops)};
+    }
+
+    auto operandShape = operandType.getShape();
+    SmallVector<AffineExpr, 4> dimExprs;
+    dimExprs.reserve(nloops);
+
+    if (broadcastOp.getBroadcastDimensions()) {
+      for (const auto& broadcastDim :
+           enumerate(broadcastOp.getBroadcastDimensions().getValues<APInt>())) {
+        int size = broadcastDim.value().getSExtValue();
+        bool expansionNeeded = operandShape[broadcastDim.index()] == 1 &&
+                               resultType.getShape()[size] != 1;
+        dimExprs.push_back(expansionNeeded ? b->getAffineConstantExpr(0)
+                                           : b->getAffineDimExpr(size));
+      }
+    }
+    return {
+        AffineMap::get(nloops, /*symbolCount=*/0, dimExprs, b->getContext()),
+        b->getMultiDimIdentityMap(nloops)};
+  }
+};
+
+Value collapseExpandingDims(PatternRewriter& rewriter, Location loc,
+                            Value operand, SmallVector<int64_t>& dimensions,
+                            llvm::function_ref<bool(int64_t)> isExpandingDim) {
+  auto operandTy = operand.getType().cast<RankedTensorType>();
+
+  SmallVector<ReassociationIndices> reassociationMap;
+  ReassociationIndices currentIndices;
+
+  ArrayRef<int64_t> operandShape = operandTy.getShape();
+  SmallVector<int64_t> newOperandShape;
+  SmallVector<int64_t> newDimensions;
+
+  for (const auto& [idx, dim] : llvm::enumerate(dimensions)) {
+    currentIndices.push_back(idx);
+
+    if (!isExpandingDim(idx)) {
+      reassociationMap.push_back(currentIndices);
+      currentIndices.clear();
+      newOperandShape.push_back(operandShape[idx]);
+      newDimensions.push_back(dim);
+    }
+  }
+
+  if (!reassociationMap.empty())
+    reassociationMap.back().insert(reassociationMap.back().end(),
+                                   currentIndices.begin(),
+                                   currentIndices.end());
+
+  if (dimensions.size() != newDimensions.size()) {
+    dimensions = newDimensions;
+
+    auto newOperandType =
+        RankedTensorType::get(newOperandShape, operandTy.getElementType());
+    operand = rewriter.create<tensor::CollapseShapeOp>(
+        loc, newOperandType, operand, reassociationMap);
+  }
+  return operand;
+}
+
+// Insert linalg.transpose if broadcasted dimensions are not in sorded order.
+// linalg.broadcast does not support implicit transpose, so the input needs to
+// be explicitely transposed.
+Value transposeBroadcastOperand(PatternRewriter& rewriter, Location loc,
+                                Value operand,
+                                SmallVector<int64_t>& dimensions) {
+  // Do not insert `transpose` is dimensions are already sorted.
+  if (llvm::is_sorted(dimensions)) return operand;
+
+  SmallVector<int64_t> permutation =
+      llvm::to_vector(llvm::seq<int64_t>(0, dimensions.size()));
+  llvm::sort(permutation, [&](int64_t lhs, int64_t rhs) {
+    return dimensions[lhs] < dimensions[rhs];
+  });
+
+  auto operandTy = operand.getType().cast<ShapedType>();
+  ArrayRef<int64_t> operandShape = operandTy.getShape();
+  SmallVector<int64_t> transposedOperandShape, transposedDimensions;
+
+  for (int64_t index : permutation) {
+    transposedOperandShape.push_back(operandShape[index]);
+    transposedDimensions.push_back(dimensions[index]);
+  }
+  dimensions = transposedDimensions;
+
+  return rewriter.create<stablehlo::TransposeOp>(
+      loc,
+      RankedTensorType::get(transposedOperandShape, operandTy.getElementType()),
+      operand, rewriter.getI64VectorAttr(permutation));
+}
+
+class BroadcastInDimOpToBroadcastConverter
+    : public OpConversionPattern<stablehlo::BroadcastInDimOp> {
+ public:
+  using OpConversionPattern<stablehlo::BroadcastInDimOp>::OpConversionPattern;
+
+  LogicalResult matchAndRewrite(
+      stablehlo::BroadcastInDimOp op, OpAdaptor adaptor,
+      ConversionPatternRewriter& rewriter) const override {
+    Location loc = op.getLoc();
+
+    SmallVector<int64_t> broadcastDimensions =
+        llvm::to_vector(op.getBroadcastDimensions().getValues<int64_t>());
+
+    Value operand = adaptor.getOperand();
+    auto operandTy = operand.getType().cast<ShapedType>();
+    auto resultTy = typeConverter->convertType(op.getType()).cast<ShapedType>();
+
+    ArrayRef<int64_t> operandShape = operandTy.getShape();
+    ArrayRef<int64_t> resultShape = resultTy.getShape();
+
+    operand = collapseExpandingDims(
+        rewriter, loc, operand, broadcastDimensions, [&](int64_t i) {
+          return operandShape[i] == 1 &&
+                 resultShape[broadcastDimensions[i]] != 1;
+        });
+    operand =
+        transposeBroadcastOperand(rewriter, loc, operand, broadcastDimensions);
+
+    Value emptyTensor =
+        getEmptyTensorFor(rewriter, loc, resultTy, op, adaptor.getOperands());
+
+    SmallVector<int64_t> addedDimensions;
+    for (int64_t dim : llvm::seq<int64_t>(0, resultTy.getRank())) {
+      if (!llvm::is_contained(broadcastDimensions, dim))
+        addedDimensions.push_back(dim);
+    }
+
+    rewriter.replaceOpWithNewOp<linalg::BroadcastOp>(
+        op, operand, emptyTensor, addedDimensions,
+        linalg::getPrunedAttributeList(op));
+    return success();
+  }
+};
+
+// If the input has a static shape we know exactly when the broadcast must
+// expand (the dimension is 1, which also trivially expands to 1) or will never
+// expand (the dimension is not 1). We can also source the information from the
+// optionally provided attrbibutes on statically known broadcasting behavior.
+// This means we can lower the broadcast just as we would lower a fully static
+// broadcast and go directly to `linalg.generic`.
+
+// This also covers the important case of broadcasting a scalar. Ideally the
+// pattern (`mhlo.constant` -> `mhlo.dynamic_broadcast_in_dim`) should be
+// converted to a tensor dialect op similar to TF's `ConstantLikeOp`.
+class HloDynamicBroadcastInDimConverter
+    : public OpConversionPattern<stablehlo::DynamicBroadcastInDimOp> {
+ public:
+  using OpConversionPattern<
+      stablehlo::DynamicBroadcastInDimOp>::OpConversionPattern;
+
+  LogicalResult matchAndRewrite(
+      stablehlo::DynamicBroadcastInDimOp op, OpAdaptor adaptor,
+      ConversionPatternRewriter& rewriter) const final {
+    Value operand = adaptor.getOperand();
+    auto operandType = operand.getType().dyn_cast<RankedTensorType>();
+    if (!operandType) return failure();
+    auto resultType =
+        typeConverter->convertType(op.getType()).dyn_cast<RankedTensorType>();
+    if (!resultType) return failure();
+
+    // Determine dimension expressions based on whether the dimension is
+    // expanding (0) or non-expanding (identity), and fail if we cannot decide
+    // this.
+    SmallVector<AffineExpr> dimExprs(operandType.getRank(), nullptr);
+
+    // Use static type info.
+    auto bcastDims = llvm::to_vector(
+        llvm::map_range(op.getBroadcastDimensions(), [](const APInt& d) {
+          return static_cast<int64_t>(d.getLimitedValue());
+        }));
+    for (const auto& it : llvm::enumerate(operandType.getShape())) {
+      if (ShapedType::isDynamic(it.value())) continue;
+      bool isExpanding = it.value() == 1;
+      dimExprs[it.index()] =
+          isExpanding ? rewriter.getAffineConstantExpr(0)
+                      : rewriter.getAffineDimExpr(bcastDims[it.index()]);
+    }
+
+    // Use annotated expansion behavior, if available.
+    if (op.getKnownExpandingDimensions()) {
+      for (const auto& it :
+           op.getKnownExpandingDimensions()->getValues<APInt>()) {
+        auto i = it.getLimitedValue();
+        dimExprs[i] = rewriter.getAffineConstantExpr(0);
+      }
+    }
+    if (op.getKnownNonexpandingDimensions()) {
+      for (const auto& it :
+           op.getKnownNonexpandingDimensions()->getValues<APInt>()) {
+        auto i = it.getLimitedValue();
+        dimExprs[i] = rewriter.getAffineDimExpr(bcastDims[i]);
+      }
+    }
+
+    // Fail if unknown expansion behavior remains.
+    if (!llvm::all_of(dimExprs, [](AffineExpr expr) { return expr; }))
+      return failure();
+
+    // Materialize `linalg.generic` op.
+    Location loc = op.getLoc();
+    int64_t nloops = resultType.getRank();
+    Value emptyTensor =
+        getEmptyTensorFor(rewriter, loc, resultType, op, adaptor.getOperands());
+    rewriter.replaceOpWithNewOp<linalg::GenericOp>(
+        op, TypeRange{emptyTensor.getType()}, ValueRange{operand},
+        /*outputBuffers=*/ValueRange{emptyTensor},
+        llvm::ArrayRef({AffineMap::get(/*dimCount=*/nloops, /*symbolCount=*/0,
+                                       dimExprs, rewriter.getContext()),
+                        rewriter.getMultiDimIdentityMap(nloops)}),
+        getNParallelLoopsAttrs(nloops),
+        [&](OpBuilder& nestedBuilder, Location /*nested_loc*/,
+            ValueRange args) {
+          nestedBuilder.create<linalg::YieldOp>(loc, *args.begin());
+        },
+        linalg::getPrunedAttributeList(op));
+    return success();
+  }
+};
+
+class DynamicBroadcastInDimOpToBroadcastConverter
+    : public OpConversionPattern<stablehlo::DynamicBroadcastInDimOp> {
+ public:
+  using OpConversionPattern<
+      stablehlo::DynamicBroadcastInDimOp>::OpConversionPattern;
+
+  LogicalResult matchAndRewrite(
+      stablehlo::DynamicBroadcastInDimOp op, OpAdaptor adaptor,
+      ConversionPatternRewriter& rewriter) const final {
+    Location loc = op.getLoc();
+
+    Value operand = adaptor.getOperand();
+    auto operandTy = operand.getType().dyn_cast<RankedTensorType>();
+    if (!operandTy) return failure();
+    auto resultTy =
+        typeConverter->convertType(op.getType()).dyn_cast<RankedTensorType>();
+    if (!resultTy) return failure();
+
+    SmallVector<int64_t> broadcastDimensions =
+        llvm::to_vector(op.getBroadcastDimensions().getValues<int64_t>());
+
+    SmallVector<std::optional<bool>> expansionBehavior(
+        broadcastDimensions.size());
+
+    // Use static type info.
+    for (const auto& [idx, dim] : llvm::enumerate(operandTy.getShape())) {
+      if (ShapedType::isDynamic(dim)) continue;
+      expansionBehavior[idx] = (dim == 1);
+    }
+
+    // Use annotated expansion behavior, if available.
+    if (op.getKnownExpandingDimensions()) {
+      for (const auto& it :
+           op.getKnownExpandingDimensions()->getValues<int64_t>()) {
+        expansionBehavior[it] = true;
+      }
+    }
+    if (op.getKnownNonexpandingDimensions()) {
+      for (const auto& it :
+           op.getKnownNonexpandingDimensions()->getValues<int64_t>()) {
+        expansionBehavior[it] = false;
+      }
+    }
+
+    // Fail if unknown expansion behavior remains.
+    if (llvm::any_of(expansionBehavior, [](auto v) { return !v.has_value(); }))
+      return failure();
+
+    auto isExpandingDim = [&](int64_t i) {
+      return expansionBehavior[i].value();
+    };
+
+    // Use attribute information to insert 1s into operand type.
+    operand = getBroadcastOperand(rewriter, loc, operand, isExpandingDim);
+
+    auto broadcastResultTy = getBroadcastResultType(
+        operand, resultTy, broadcastDimensions, isExpandingDim);
+
+    operand = collapseExpandingDims(rewriter, loc, operand, broadcastDimensions,
+                                    isExpandingDim);
+    operand =
+        transposeBroadcastOperand(rewriter, loc, operand, broadcastDimensions);
+
+    Value emptyTensor = getEmptyTensorFor(rewriter, loc, broadcastResultTy, op,
+                                          adaptor.getOperands());
+
+    SmallVector<int64_t> addedDimensions;
+    for (int64_t dim : llvm::seq<int64_t>(0, resultTy.getRank())) {
+      if (!llvm::is_contained(broadcastDimensions, dim))
+        addedDimensions.push_back(dim);
+    }
+
+    Value result = rewriter
+                       .create<linalg::BroadcastOp>(
+                           loc, operand, emptyTensor, addedDimensions,
+                           linalg::getPrunedAttributeList(op))
+                       .getResults()[0];
+
+    if (resultTy != broadcastResultTy) {
+      result = rewriter.create<tensor::CastOp>(loc, resultTy, result);
+    }
+
+    rewriter.replaceOp(op, result);
+    return success();
+  }
+
+ private:
+  static Value getBroadcastOperand(
+      PatternRewriter& rewriter, Location loc, Value operand,
+      llvm::function_ref<bool(int64_t)> isExpandingDim) {
+    auto operandTy = operand.getType().dyn_cast<RankedTensorType>();
+
+    SmallVector<int64_t> updatedOperandShape =
+        llvm::to_vector(operandTy.getShape());
+    for (size_t i = 0; i < updatedOperandShape.size(); ++i) {
+      if (isExpandingDim(i)) updatedOperandShape[i] = 1;
+    }
+
+    auto updatedOperandTy =
+        RankedTensorType::get(updatedOperandShape, operandTy.getElementType());
+
+    if (updatedOperandTy != operandTy) {
+      operand = rewriter.create<tensor::CastOp>(loc, updatedOperandTy, operand);
+    }
+
+    return operand;
+  }
+
+  static ShapedType getBroadcastResultType(
+      Value operand, RankedTensorType resultTy, ArrayRef<int64_t> dimensions,
+      llvm::function_ref<bool(int64_t)> isExpandingDim) {
+    auto operandShape = operand.getType().cast<RankedTensorType>().getShape();
+    auto broadcastResultShape = llvm::to_vector(resultTy.getShape());
+
+    for (auto [operandIndex, resultIndex] : llvm::enumerate(dimensions)) {
+      if (isExpandingDim(operandIndex)) continue;
+      broadcastResultShape[resultIndex] = operandShape[operandIndex];
+    }
+
+    return RankedTensorType::get(broadcastResultShape,
+                                 resultTy.getElementType());
+  }
+};
+
+template <typename OpTy>
+class TransposeConverter
+    : public DataMovementOpConverter<TransposeConverter<OpTy>, OpTy> {
+ public:
+  using DataMovementOpConverter<TransposeConverter<OpTy>,
+                                OpTy>::DataMovementOpConverter;
+  static SmallVector<AffineMap, 2> getIndexingMaps(OpTy op, Builder* b) {
+    auto resultType = getHloOpResultType(op).template cast<ShapedType>();
+    auto nloops = resultType.getRank();
+    SmallVector<AffineExpr, 2> inputExprs;
+    inputExprs.resize(resultType.getRank());
+    for (const auto& permutation : llvm::enumerate(op.getPermutation())) {
+      inputExprs[permutation.value().getZExtValue()] =
+          b->getAffineDimExpr(permutation.index());
+    }
+    return {
+        AffineMap::get(nloops, /*symbolCount=*/0, inputExprs, b->getContext()),
+        b->getMultiDimIdentityMap(nloops)};
+  }
+};
+
+class TransposeOpToTransposeConverter
+    : public OpConversionPattern<stablehlo::TransposeOp> {
+  using OpConversionPattern<stablehlo::TransposeOp>::OpConversionPattern;
+
+  LogicalResult matchAndRewrite(
+      stablehlo::TransposeOp op, OpAdaptor adaptor,
+      ConversionPatternRewriter& rewriter) const override {
+    auto resultTy = typeConverter->convertType(op.getType()).cast<ShapedType>();
+
+    auto loc = op.getLoc();
+    Value emptyTensor =
+        getEmptyTensorFor(rewriter, loc, resultTy, op, adaptor.getOperands());
+
+    auto permutation = rewriter.getDenseI64ArrayAttr(
+        llvm::to_vector(op.getPermutation().getValues<int64_t>()));
+
+    rewriter.replaceOpWithNewOp<linalg::TransposeOp>(
+        op, adaptor.getOperand(), emptyTensor, permutation,
+        linalg::getPrunedAttributeList(op));
+    return success();
+  }
+};
+
+class BitcastConvertConverter
+    : public OpConversionPattern<stablehlo::BitcastConvertOp> {
+  using OpConversionPattern::OpConversionPattern;
+
+  LogicalResult matchAndRewrite(
+      stablehlo::BitcastConvertOp op, OpAdaptor adaptor,
+      ConversionPatternRewriter& rewriter) const final {
+    if (!verifyHloOpBufferOrTensorSemantics(op)) return failure();
+
+    auto inputType = adaptor.getOperand().getType().cast<RankedTensorType>();
+    auto outputType =
+        typeConverter->convertType(op.getType()).cast<RankedTensorType>();
+    auto loc = op.getLoc();
+
+    // Fallback to pointwise conversion if the tensor dimensions are not
+    // changing.
+    if (inputType.getRank() == outputType.getRank()) {
+      return failure();
+    }
+
+    auto inputBitWidth = inputType.getElementType().getIntOrFloatBitWidth();
+    auto outputBitWidth = outputType.getElementType().getIntOrFloatBitWidth();
+
+    auto maxRank = std::max(inputType.getRank(), outputType.getRank());
+    auto identityMap =
+        AffineMap::getMultiDimIdentityMap(maxRank, rewriter.getContext());
+    AffineMap indexingMaps[] = {
+        AffineMap::get(
+            /*dimCount=*/maxRank, /*symbolCount=*/0,
+            identityMap.getResults().take_front(inputType.getRank()),
+            rewriter.getContext()),
+        AffineMap::get(
+            /*dimCount=*/maxRank, /*symbolCount=*/0,
+            identityMap.getResults().take_front(outputType.getRank()),
+            rewriter.getContext())};
+
+    Value output =
+        getEmptyTensorFor(rewriter, loc, outputType, op, adaptor.getOperands());
+    bool isExpansion = inputBitWidth > outputBitWidth;
+    bool isContraction = inputBitWidth < outputBitWidth;
+    // When combining values we start with a 0 and merge bits into it.
+    if (isContraction) {
+      output = fillTensorWithZeros(rewriter, loc, output);
+    }
+
+    rewriter.replaceOpWithNewOp<linalg::GenericOp>(
+        op, outputType, adaptor.getOperand(), output, indexingMaps,
+        getParallelAndReductionIterators(maxRank, isContraction ? 1 : 0),
+        [&](OpBuilder& nestedBuilder, Location nestedLoc, ValueRange args) {
+          auto inIntType = nestedBuilder.getIntegerType(inputBitWidth);
+          auto outIntType = nestedBuilder.getIntegerType(outputBitWidth);
+          Value innerResult = args.front();
+          if (isExpansion) {
+            // Expand a big value into multiple small values with shifts.
+            auto iotaIndex =
+                nestedBuilder.create<linalg::IndexOp>(nestedLoc, maxRank - 1);
+            auto iota = nestedBuilder.create<arith::IndexCastOp>(
+                nestedLoc, inIntType, iotaIndex);
+
+            auto width = nestedBuilder.create<arith::ConstantOp>(
+                nestedLoc,
+                nestedBuilder.getIntegerAttr(inIntType, outputBitWidth));
+            auto shiftWidth =
+                nestedBuilder.create<arith::MulIOp>(nestedLoc, iota, width);
+            Value inputCasted = nestedBuilder.create<arith::BitcastOp>(
+                nestedLoc, inIntType, args.front());
+            Value shifted = nestedBuilder.create<arith::ShRUIOp>(
+                nestedLoc, inputCasted, shiftWidth);
+            innerResult = nestedBuilder.create<arith::TruncIOp>(
+                nestedLoc, outIntType, shifted);
+          } else if (isContraction) {
+            // Combine multiple small values into one big value.
+            auto iotaIndex =
+                nestedBuilder.create<linalg::IndexOp>(nestedLoc, maxRank - 1);
+            auto iota = nestedBuilder.create<arith::IndexCastOp>(
+                nestedLoc, outIntType, iotaIndex);
+
+            auto width = nestedBuilder.create<arith::ConstantOp>(
+                nestedLoc,
+                nestedBuilder.getIntegerAttr(outIntType, inputBitWidth));
+            auto shiftWidth =
+                nestedBuilder.create<arith::MulIOp>(nestedLoc, iota, width);
+            Value inputCasted = nestedBuilder.create<arith::BitcastOp>(
+                nestedLoc, inIntType, args.front());
+            Value inputExt = nestedBuilder.create<arith::ExtUIOp>(
+                nestedLoc, outIntType, inputCasted);
+            Value shifted = nestedBuilder.create<arith::ShLIOp>(
+                nestedLoc, inputExt, shiftWidth);
+            Value accumulatorCasted = nestedBuilder.create<arith::BitcastOp>(
+                nestedLoc, outIntType, args.back());
+            innerResult = nestedBuilder.create<arith::OrIOp>(
+                nestedLoc, outIntType, shifted, accumulatorCasted);
+          }
+          innerResult = nestedBuilder.create<arith::BitcastOp>(
+              nestedLoc, outputType.getElementType(), innerResult);
+          nestedBuilder.create<linalg::YieldOp>(nestedLoc, innerResult);
+        },
+        linalg::getPrunedAttributeList(op));
+    return success();
+  }
+};
+
+// Lowers mhlo.RealDynamicSliceOp to tensor.extract_slice and other
+// arith/tensor dialect ops.
+class RealDynamicSliceConverter
+    : public OpConversionPattern<stablehlo::RealDynamicSliceOp> {
+ public:
+  using OpConversionPattern<stablehlo::RealDynamicSliceOp>::OpConversionPattern;
+
+  // Computes size of a slice as
+  //   size = ceil((limit - start)/stride)
+  static Value computeSize(Location loc, Value start, Value limit, Value stride,
+                           ConversionPatternRewriter& b) {
+    Value delta = b.create<arith::SubIOp>(loc, limit, start);
+    Value ret = b.create<arith::CeilDivUIOp>(loc, delta, stride);
+    if (ret.getType().isIndex()) return ret;
+    return b.create<arith::IndexCastOp>(loc, b.getIndexType(), ret);
+  }
+
+  LogicalResult matchAndRewrite(
+      stablehlo::RealDynamicSliceOp realDynamicSliceOp, OpAdaptor adaptor,
+      ConversionPatternRewriter& rewriter) const final {
+    Location loc = realDynamicSliceOp.getLoc();
+    auto argType = adaptor.getOperand().getType().dyn_cast<ShapedType>();
+    if (!argType || !argType.hasRank()) {
+      return rewriter.notifyMatchFailure(realDynamicSliceOp,
+                                         "require known-rank args");
+    }
+
+    Type dimElementType = getElementTypeOrSelf(adaptor.getStartIndices());
+    if (getElementTypeOrSelf(adaptor.getLimitIndices()) != dimElementType ||
+        getElementTypeOrSelf(adaptor.getStrides()) != dimElementType) {
+      return rewriter.notifyMatchFailure(
+          realDynamicSliceOp,
+          "requires same element type for all dimension specification");
+    }
+    Type arithType =
+        dimElementType.isIndex() ? rewriter.getI64Type() : dimElementType;
+    Type indexType = rewriter.getIndexType();
+
+    auto resultType =
+        this->typeConverter->convertType(realDynamicSliceOp.getType())
+            .cast<RankedTensorType>();
+    Value zero = rewriter.create<arith::ConstantIndexOp>(loc, 0);
+    SmallVector<OpFoldResult, 4> offsets, sizes, strides;
+    SmallVector<Type, 3> clampType(3, arithType);
+    for (auto i : llvm::seq<unsigned>(0, argType.getRank())) {
+      Value dim = rewriter.create<arith::ConstantIndexOp>(loc, i);
+      Value start = rewriter.create<tensor::ExtractOp>(
+          loc, adaptor.getStartIndices(), dim);
+      Value limit = rewriter.create<tensor::ExtractOp>(
+          loc, adaptor.getLimitIndices(), dim);
+      Value stride =
+          rewriter.create<tensor::ExtractOp>(loc, adaptor.getStrides(), dim);
+
+      // Compute i-th dimension size of the result : size[i].
+      // If the i-th dimension of the result type is known, we go ahead with it
+      // else we compute it using limit, start and stride values.
+      int64_t resultDimSize = resultType.getDimSize(i);
+      Value size =
+          ShapedType::isDynamic(resultDimSize)
+              ? computeSize(loc, start, limit, stride, rewriter)
+              : rewriter.create<arith::ConstantIndexOp>(loc, resultDimSize);
+
+      // We can now convert start to index.
+      if (!start.getType().isIndex())
+        start = rewriter.create<arith::IndexCastOp>(
+            loc, rewriter.getIndexType(), start);
+
+      // Fetch i-th dimension size of the operand and calculate upper bound as
+      //   ub = operand_dim[i] - size[i]
+      Value operandDimSize =
+          rewriter.createOrFold<tensor::DimOp>(loc, adaptor.getOperand(), dim);
+      Value upperBound =
+          rewriter.createOrFold<arith::SubIOp>(loc, operandDimSize, size);
+
+      // We clamp the start_index to keep it bounded as
+      //   0 <= start_index[i] <= ub
+      // Clamp does not support index type, so cast to integer type.
+      start = rewriter.create<arith::MaxSIOp>(loc, start, zero);
+      start = rewriter.create<arith::MinSIOp>(loc, start, upperBound);
+
+      offsets.push_back(start);
+      if (ShapedType::isDynamic(resultDimSize))
+        sizes.push_back(size);
+      else
+        sizes.push_back(IntegerAttr::get(indexType, resultDimSize));
+
+      if (!stride.getType().isIndex())
+        stride =
+            rewriter.createOrFold<arith::IndexCastOp>(loc, indexType, stride);
+      strides.push_back(stride);
+    }
+
+    rewriter.replaceOpWithNewOp<tensor::ExtractSliceOp>(
+        realDynamicSliceOp, resultType, adaptor.getOperand(), offsets, sizes,
+        strides);
+    return success();
+  }
+};
+
+// Converts reshape ops that can be proven to be either a collapse of dimensions
+// or expansion of dimensions of the operand.
+class ReshapeOpConverter : public OpConversionPattern<stablehlo::ReshapeOp> {
+ public:
+  using OpConversionPattern::OpConversionPattern;
+
+  LogicalResult matchAndRewrite(
+      stablehlo::ReshapeOp reshapeOp, stablehlo::ReshapeOp::Adaptor adaptor,
+      ConversionPatternRewriter& rewriter) const final {
+    if (!verifyHloOpBufferOrTensorSemantics(reshapeOp)) return failure();
+    auto operand = adaptor.getOperand();
+    auto operandType = operand.getType().cast<ShapedType>();
+    auto elemType = operandType.getElementType();
+    auto resultType = reshapeOp.getType().cast<ShapedType>();
+
+    if (!resultType.hasStaticShape()) return failure();
+
+    // If any of the output dimensions is 0, the tensor has no elements. In that
+    // case, we can just replace the reshape with an empty op.
+    if (llvm::is_contained(resultType.getShape(), 0)) {
+      rewriter.replaceOpWithNewOp<tensor::EmptyOp>(
+          reshapeOp, resultType.getShape(), elemType);
+      return success();
+    }
+
+    resultType = typeConverter->convertType(resultType).cast<ShapedType>();
+
+    // Special case where the result is a scalar.
+    if (resultType.getRank() == 0 && !operandType.hasStaticShape()) {
+      // This means all dimensions of the operand need to be 1. We add a cast to
+      // cast the dynamic dimensions to 1.
+      auto staticType = RankedTensorType::get(
+          llvm::SmallVector<int64_t>(operandType.getRank(), 1), elemType);
+      operand = rewriter.create<tensor::CastOp>(reshapeOp.getLoc(), staticType,
+                                                operand);
+      rewriter.replaceOpWithNewOp<tensor::CollapseShapeOp>(
+          reshapeOp, resultType, operand, ArrayRef<ReassociationIndices>{});
+      return success();
+    }
+
+    // Compute the reassociation maps for the linalg operation. This will
+    // succeed if the reshape can be done with a single expand_shape or
+    // collapse_shape.
+    if (std::optional<SmallVector<ReassociationIndices>> reassociationMap =
+            getReassociationIndicesForReshape(operandType, resultType)) {
+      if (resultType.getRank() < operandType.getRank()) {
+        // We have found a working reassociation map. If the operand is dynamic,
+        // we first need to cast all unknown dimensions in the input that get
+        // collapsed to a static-sized dimension in the output, to 1.
+        SmallVector<int64_t> shape(operandType.getShape().begin(),
+                                   operandType.getShape().end());
+        for (const auto& map : llvm::enumerate(*reassociationMap)) {
+          // If the result dim is dynamic, we do not mind dynamic entries in the
+          // source.
+          if (resultType.isDynamicDim(map.index())) continue;
+          for (auto targetDim : map.value()) {
+            if (shape[targetDim] == ShapedType::kDynamic) shape[targetDim] = 1;
+          }
+        }
+        // Insert a cast if types are not the same (ignoring sparse encoding).
+        auto enc = sparse_tensor::getSparseTensorEncoding(operandType);
+        auto newOperandType = RankedTensorType::get(shape, elemType, enc);
+        if (newOperandType != operandType) {
+          operand = rewriter.create<tensor::CastOp>(reshapeOp.getLoc(),
+                                                    newOperandType, operand);
+        }
+        // Generate collapse operation.
+        rewriter.replaceOpWithNewOp<tensor::CollapseShapeOp>(
+            reshapeOp, resultType, operand, *reassociationMap);
+      } else {
+        // Generate expand operation.
+        rewriter.replaceOpWithNewOp<tensor::ExpandShapeOp>(
+            reshapeOp, resultType, operand, *reassociationMap);
+      }
+      return success();
+    }
+
+    Value collapsedOp = operand;
+    Location loc = reshapeOp.getLoc();
+    auto getIdentityExprs = [&rewriter](int64_t n) {
+      SmallVector<AffineExpr, 4> exprs;
+      for (int i = 0; i < n; ++i) exprs.push_back(rewriter.getAffineDimExpr(i));
+      return exprs;
+    };
+    // Otherwise, we need to first reduce all source dimensions into one and
+    // then expand to the destination dimensions. If there is only a single
+    // source dimension, the reduce step can be skipped. TensorCollapseShape
+    // expects a different rank of operand and result.
+    if (operandType.getRank() != 1) {
+      SmallVector<ReassociationExprs, 4> collapsingMap = {
+          // Use operand_type here because we need to collapse all operands
+          // dimensions.
+          getIdentityExprs(operandType.getRank())};
+
+      collapsedOp =
+          rewriter.create<tensor::CollapseShapeOp>(loc, operand, collapsingMap);
+    }
+    // Cast to a known static type if the input has dynamic dimensions.
+    int64_t totalElems = resultType.getNumElements();
+    auto collapsedType = RankedTensorType::get({totalElems}, elemType);
+    collapsedOp =
+        rewriter.create<tensor::CastOp>(loc, collapsedType, collapsedOp);
+    if (resultType.getRank() == 1) {
+      rewriter.replaceOp(reshapeOp, collapsedOp);
+    } else {
+      SmallVector<ReassociationExprs, 4> expandingMap = {
+          // Use resultType here because we need to expand to all result
+          // dimensions.
+          getIdentityExprs(resultType.getRank())};
+      rewriter.replaceOpWithNewOp<tensor::ExpandShapeOp>(
+          reshapeOp, resultType, collapsedOp, expandingMap);
+    }
+    return success();
+  }
+};
+
+template <typename OpTy>
+class IotaConverter : public OpConversionPattern<OpTy> {
+ public:
+  using OpConversionPattern<OpTy>::OpConversionPattern;
+
+  LogicalResult matchAndRewrite(
+      OpTy iotaOp, typename OpTy::Adaptor adaptor,
+      ConversionPatternRewriter& rewriter) const final {
+    ShapedType resultShapedType = getHloOpResultType(iotaOp);
+    if (!resultShapedType) return failure();
+    resultShapedType = this->typeConverter->convertType(resultShapedType)
+                           .template dyn_cast<ShapedType>();
+
+    Type resultElementType = resultShapedType.getElementType();
+
+    // Construct the indexing maps needed for linalg.generic ops.
+    unsigned nloops = resultShapedType.getRank();
+
+    Location loc = iotaOp.getLoc();
+    auto linalgOp = rewriter.create<linalg::GenericOp>(
+        loc,
+        /*resultTensorTypes=*/
+        ArrayRef<Type>{resultShapedType},
+        /*inputs=*/ValueRange{},
+        /*outputBuffers=*/
+
+        ValueRange{getEmptyTensorFor(rewriter, loc, resultShapedType, iotaOp,
+                                     adaptor.getOperands())},
+        llvm::ArrayRef(rewriter.getMultiDimIdentityMap(nloops)),
+        getNParallelLoopsAttrs(nloops),
+        [&](OpBuilder& nestedBuilder, Location nestedLoc, ValueRange /*args*/) {
+          Value indexOp = nestedBuilder.create<linalg::IndexOp>(
+              nestedLoc, iotaOp.getIotaDimension());
+          Type unwrappedResultElementType = resultElementType;
+          if (auto complexType =
+                  unwrappedResultElementType.dyn_cast<ComplexType>())
+            unwrappedResultElementType = complexType.getElementType();
+          Value castOp = nestedBuilder.create<arith::IndexCastOp>(
+              nestedLoc,
+              nestedBuilder.getIntegerType(
+                  unwrappedResultElementType.getIntOrFloatBitWidth()),
+              indexOp);
+          castOp = stablehlo::StableHloOpToStdScalarOp::mapOpOfType<
+              stablehlo::ConvertOp>(nestedLoc, resultElementType,
+                                    castOp.getType(), {castOp}, &nestedBuilder);
+          nestedBuilder.create<linalg::YieldOp>(nestedLoc, castOp);
+        },
+        linalg::getPrunedAttributeList(iotaOp));
+    rewriter.replaceOp(iotaOp, linalgOp.getResultTensors());
+    return success();
+  }
+};
+
+template <typename OpTy>
+class IotaToMapConverter : public OpConversionPattern<OpTy> {
+ public:
+  using OpConversionPattern<OpTy>::OpConversionPattern;
+
+  LogicalResult matchAndRewrite(
+      OpTy iotaOp, typename OpTy::Adaptor adaptor,
+      ConversionPatternRewriter& rewriter) const final {
+    ShapedType resultTy = getHloOpResultType(iotaOp);
+    if (!resultTy) return failure();
+    resultTy = this->typeConverter->convertType(resultTy)
+                   .template dyn_cast<ShapedType>();
+
+    Location loc = iotaOp.getLoc();
+    Value empty = getEmptyTensorFor(rewriter, loc, resultTy, iotaOp,
+                                    adaptor.getOperands());
+
+    auto linalgOp = rewriter.create<linalg::MapOp>(
+        loc, ValueRange{}, empty,
+        [&](OpBuilder& nestedBuilder, Location nestedLoc, ValueRange /*args*/) {
+          Value index = nestedBuilder.create<linalg::IndexOp>(
+              nestedLoc, iotaOp.getIotaDimension());
+          index = nestedBuilder.create<arith::IndexCastOp>(
+              nestedLoc, nestedBuilder.getI64Type(), index);
+          Value result = stablehlo::StableHloOpToStdScalarOp::mapOpOfType<
+              stablehlo::ConvertOp>(nestedLoc, resultTy.getElementType(),
+                                    index.getType(), {ValueRange{index}},
+                                    &nestedBuilder);
+          nestedBuilder.create<linalg::YieldOp>(nestedLoc, ValueRange{result});
+        },
+        linalg::getPrunedAttributeList(iotaOp));
+    rewriter.replaceOp(iotaOp, linalgOp.getResult());
+    return success();
+  }
+};
+
+/// Converts stablehlo.concatenate operation to a linalg.generic op.
+struct ConcatenateConverter
+    : public OpConversionPattern<stablehlo::ConcatenateOp> {
+  using OpConversionPattern::OpConversionPattern;
+
+  LogicalResult matchAndRewrite(
+      stablehlo::ConcatenateOp op, OpAdaptor adaptor,
+      ConversionPatternRewriter& rewriter) const override {
+    // Shortcut the one-operand case, simplifies code below.
+    if (adaptor.getOperands().size() == 1) {
+      rewriter.replaceOp(op, adaptor.getOperands()[0]);
+      return success();
+    }
+
+    auto resultType = this->typeConverter->convertType(op.getResult().getType())
+                          .dyn_cast<RankedTensorType>();
+    if (!resultType) return failure();
+
+    uint64_t dim = op.getDimension();
+    Location loc = op.getLoc();
+    Value zero = rewriter.create<arith::ConstantIndexOp>(loc, 0);
+
+    // Allocate the output tensor with tensor.empty.
+    Value result =
+        getEmptyTensorFor(rewriter, loc, resultType, op, adaptor.getOperands());
+
+    // Generate a generic op to gather the elements of the concatenate. This is
+    // awkward standalone but allows fusion with other generic ops.
+    int64_t nloops = resultType.getRank();
+    rewriter.replaceOpWithNewOp<linalg::GenericOp>(
+        op,
+        /*resultTensorTypes=*/resultType,
+        /*inputs=*/ValueRange{}, /*outputBuffers=*/result,
+        llvm::ArrayRef(rewriter.getMultiDimIdentityMap(nloops)),
+        getNParallelLoopsAttrs(nloops),
+        [&](OpBuilder& nestedBuilder, Location loc, ValueRange) {
+          OpBuilder b = nestedBuilder;
+          Value concatDimSize = zero;
+          Value result;
+
+          SmallVector<Value, 4> extractIndices;
+          extractIndices.reserve(nloops);
+          for (int64_t i = 0; i < nloops; i++) {
+            extractIndices.push_back(b.create<linalg::IndexOp>(loc, i));
+          }
+
+          Value indexOp = b.create<linalg::IndexOp>(loc, dim);
+          for (const auto& it : llvm::enumerate(adaptor.getOperands())) {
+            Value arg = it.value();
+            Value newConcatDimSize;
+            scf::IfOp ifOp;
+            if (it.index() != (adaptor.getOperands().size() - 1)) {
+              // Calculate how far along we have iterated along the concatenate
+              // dimension. That way we can tell which input to select.
+              newConcatDimSize = b.create<arith::AddIOp>(
+                  loc, concatDimSize, b.create<tensor::DimOp>(loc, arg, dim));
+              Value cmp = b.create<arith::CmpIOp>(loc, rewriter.getI1Type(),
+                                                  arith::CmpIPredicate::ult,
+                                                  indexOp, newConcatDimSize);
+              ifOp = b.create<scf::IfOp>(loc, resultType.getElementType(), cmp,
+                                         true);
+              if (result) {
+                b.create<scf::YieldOp>(loc, ifOp->getResults()[0]);
+              } else {
+                result = ifOp->getResults()[0];
+              }
+
+              b = ifOp.getThenBodyBuilder(b.getListener());
+            }
+
+            // Now adjust the index for the concatenated dimension to fit into
+            // the selected tensor and do an extract at that position.
+            extractIndices[dim] =
+                b.create<arith::SubIOp>(loc, indexOp, concatDimSize);
+            Value extract =
+                b.create<tensor::ExtractOp>(loc, arg, extractIndices);
+            b.create<scf::YieldOp>(loc, extract);
+
+            if (ifOp) {
+              b = ifOp.getElseBodyBuilder(b.getListener());
+              concatDimSize = newConcatDimSize;
+            }
+          }
+          nestedBuilder.create<linalg::YieldOp>(loc, result);
+        },
+        linalg::getPrunedAttributeList(op));
+    return success();
+  }
+};
+
+class ConstConverterTensor : public OpConversionPattern<stablehlo::ConstantOp> {
+ public:
+  using OpConversionPattern::OpConversionPattern;
+
+  LogicalResult matchAndRewrite(
+      stablehlo::ConstantOp constOp, OpAdaptor /*adaptor*/,
+      ConversionPatternRewriter& rewriter) const final {
+    auto valueAttr = constOp.getValue().cast<DenseElementsAttr>();
+    auto type =
+        typeConverter->convertType(constOp.getType()).cast<ShapedType>();
+    if (type != constOp.getType()) {
+      // Signedness conversion.
+      valueAttr = valueAttr.mapValues(type.getElementType(),
+                                      [](const APInt& i) { return i; });
+    }
+    rewriter.replaceOpWithNewOp<arith::ConstantOp>(constOp, type, valueAttr);
+    return success();
+  }
+};
+
+// TODO(b/156787842): Support the lowering for dynamic shapes.
+class ReverseConverter
+    : public DataMovementOpConverter<ReverseConverter, stablehlo::ReverseOp> {
+ public:
+  using DataMovementOpConverter<ReverseConverter,
+                                stablehlo::ReverseOp>::DataMovementOpConverter;
+  static SmallVector<AffineMap, 2> getIndexingMaps(stablehlo::ReverseOp op,
+                                                   Builder* b) {
+    auto resultType = getHloOpResultType(op).cast<ShapedType>();
+    auto nloops = resultType.getRank();
+    SmallVector<AffineExpr, 2> inputExprs;
+    inputExprs.reserve(nloops);
+    for (int i = 0; i < nloops; ++i)
+      inputExprs.push_back(b->getAffineDimExpr(i));
+    for (auto dim : op.getDimensions()) {
+      int i = dim.getZExtValue();
+      if (resultType.isDynamicDim(i)) return {};
+      int n = resultType.getShape()[i];
+      inputExprs[i] = b->getAffineConstantExpr(n - 1) - inputExprs[i];
+    }
+    return {
+        AffineMap::get(nloops, /*symbolCount=*/0, inputExprs, b->getContext()),
+        b->getMultiDimIdentityMap(nloops)};
+  }
+};
+
+class SliceConverter : public OpConversionPattern<stablehlo::SliceOp> {
+ public:
+  using OpConversionPattern::OpConversionPattern;
+
+  LogicalResult matchAndRewrite(
+      stablehlo::SliceOp sliceOp, typename stablehlo::SliceOp::Adaptor adaptor,
+      ConversionPatternRewriter& rewriter) const final {
+    auto argType = adaptor.getOperands()[0].getType().dyn_cast<ShapedType>();
+    if (!argType || !argType.hasRank()) {
+      return rewriter.notifyMatchFailure(sliceOp, "expects known-rank args");
+    }
+
+    SmallVector<OpFoldResult, 3> offsets, sizes, strides;
+    for (int i = 0, e = argType.getRank(); i < e; ++i) {
+      auto start = sliceOp.getStartIndices().getValues<int64_t>()[i];
+      auto limit = sliceOp.getLimitIndices().getValues<int64_t>()[i];
+      auto stride = sliceOp.getStrides().getValues<int64_t>()[i];
+      offsets.push_back(rewriter.getI64IntegerAttr(start));
+      // Say that there are k elements in total, we have condition:
+      //   start + (k - 1) * strides <= limit - 1
+      // ->
+      //   k <= (limit - 1 - start + strides) / strides
+      sizes.push_back(
+          rewriter.getI64IntegerAttr((limit - 1 - start + stride) / stride));
+      strides.push_back(rewriter.getI64IntegerAttr(stride));
+    }
+    rewriter.replaceOpWithNewOp<tensor::ExtractSliceOp>(
+        sliceOp, adaptor.getOperands()[0], offsets, sizes, strides);
+    return success();
+  }
+};
+
+class DynamicSliceConverter
+    : public OpConversionPattern<stablehlo::DynamicSliceOp> {
+ public:
+  using OpConversionPattern<stablehlo::DynamicSliceOp>::OpConversionPattern;
+
+  LogicalResult matchAndRewrite(
+      stablehlo::DynamicSliceOp dynamicSliceOp, OpAdaptor adaptor,
+      ConversionPatternRewriter& rewriter) const final {
+    auto loc = dynamicSliceOp.getLoc();
+    auto argType = adaptor.getOperand().getType().dyn_cast<ShapedType>();
+    if (!argType || !argType.hasRank()) {
+      return rewriter.notifyMatchFailure(dynamicSliceOp,
+                                         "require known-rank args");
+    }
+
+    SmallVector<OpFoldResult, 3> startIndices, sizes;
+    Type originalStartIndexType =
+        dynamicSliceOp.getStartIndices().front().getType();
+    for (const auto& en : llvm::enumerate(
+             llvm::zip(adaptor.getStartIndices(),
+                       dynamicSliceOp.getSliceSizes().getValues<int64_t>()))) {
+      int64_t size = std::get<1>(en.value());
+      sizes.push_back(rewriter.getI64IntegerAttr(size));
+
+      // By mhlo.DynamicSlice definition:
+      //   `start_indices[i] = clamp(start_indices[i],
+      //       0, operand.dimension_size[i] - size_indices[i])`
+      Value startIndex = extractIndexFromTensor(
+          rewriter, loc, std::get<0>(en.value()), originalStartIndexType);
+
+      Value mn = rewriter.create<arith::ConstantIndexOp>(loc, 0);
+
+      Value mx = rewriter.createOrFold<tensor::DimOp>(loc, adaptor.getOperand(),
+                                                      en.index());
+      mx = rewriter.createOrFold<arith::SubIOp>(
+          loc, mx, rewriter.create<arith::ConstantIndexOp>(loc, size));
+
+      startIndex = rewriter.create<arith::MaxSIOp>(loc, startIndex, mn);
+      startIndex = rewriter.create<arith::MinSIOp>(loc, startIndex, mx);
+
+      startIndices.push_back(startIndex);
+    }
+
+    int64_t rank = argType.getRank();
+    SmallVector<OpFoldResult, 3> strides(rank, rewriter.getI64IntegerAttr(1));
+
+    auto resultType = this->typeConverter->convertType(dynamicSliceOp.getType())
+                          .cast<RankedTensorType>();
+
+    rewriter.replaceOpWithNewOp<tensor::ExtractSliceOp>(
+        dynamicSliceOp, resultType, adaptor.getOperand(), startIndices, sizes,
+        strides);
+    return success();
+  }
+};
+
+class DynamicUpdateSliceConverter
+    : public OpConversionPattern<stablehlo::DynamicUpdateSliceOp> {
+ public:
+  using OpConversionPattern<
+      stablehlo::DynamicUpdateSliceOp>::OpConversionPattern;
+
+  LogicalResult matchAndRewrite(
+      stablehlo::DynamicUpdateSliceOp op, OpAdaptor adaptor,
+      ConversionPatternRewriter& rewriter) const final {
+    auto loc = op.getLoc();
+    auto operandType =
+        adaptor.getOperand().getType().dyn_cast<RankedTensorType>();
+    if (!operandType || !operandType.hasStaticShape()) {
+      return rewriter.notifyMatchFailure(
+          op, "require static ranked type for operand");
+    }
+
+    auto updateType =
+        adaptor.getUpdate().getType().dyn_cast<RankedTensorType>();
+    if (!updateType || !updateType.hasStaticShape()) {
+      return rewriter.notifyMatchFailure(
+          op, "require static ranked type for operand");
+    }
+
+    // We do not have to clamp sizes because the semantic of `update`
+    // guarantees that it is always in the bounds. See
+    // https://www.tensorflow.org/xla/operation_semantics#dynamicupdateslice
+    SmallVector<OpFoldResult, 3> sizes;
+    for (auto size : updateType.getShape()) {
+      sizes.push_back(rewriter.getIndexAttr(size));
+    }
+
+    SmallVector<OpFoldResult, 3> startIndices;
+    Value zero = rewriter.create<arith::ConstantIndexOp>(loc, 0);
+    for (const auto& en : llvm::enumerate(adaptor.getStartIndices())) {
+      // By mhlo.DynamicUpdateSlice definition:
+      //   `start_indices[i] = clamp(start_indices[i],
+      //       0, operand.dimension_size[i] - update.dimension_size[i])`
+      Value startIndex =
+          extractIndexFromTensor(rewriter, loc, en.value(),
+                                 op.getStartIndices()[en.index()].getType());
+      Value ub = rewriter.create<arith::ConstantIndexOp>(
+          loc, operandType.getDimSize(en.index()) -
+                   updateType.getDimSize(en.index()));
+
+      startIndex = rewriter.create<arith::MaxSIOp>(loc, startIndex, zero);
+      startIndex = rewriter.create<arith::MinSIOp>(loc, startIndex, ub);
+      startIndices.push_back(startIndex);
+    }
+
+    int64_t rank = operandType.getRank();
+    SmallVector<OpFoldResult, 3> strides(rank, rewriter.getI64IntegerAttr(1));
+    rewriter.replaceOpWithNewOp<tensor::InsertSliceOp>(
+        op, adaptor.getUpdate(), adaptor.getOperand(), startIndices, sizes,
+        strides);
+    return success();
+  }
+};
+
+enum class DotOperationType {
+  kVectorDot = 0,
+  kMatrixVector,
+  kVectorMatrix,
+  kMatrixMatrix,
+  kUnsupported
+};
+
+DotOperationType getDotOperationType(stablehlo::DotOp dotOp) {
+  ArrayRef<int64_t> lhsShape =
+      dotOp.getLhs().getType().cast<ShapedType>().getShape();
+  ArrayRef<int64_t> rhsShape =
+      dotOp.getRhs().getType().cast<ShapedType>().getShape();
+  auto shapeMatches = [](int64_t a, int64_t b) {
+    return a == ShapedType::kDynamic || b == ShapedType::kDynamic || a == b;
+  };
+  if (lhsShape.size() == 1 && rhsShape.size() == 1 &&
+      shapeMatches(lhsShape[0], rhsShape[0])) {
+    return DotOperationType::kVectorDot;
+  }
+  if (lhsShape.size() == 2 && rhsShape.size() == 1 &&
+      shapeMatches(lhsShape[1], rhsShape[0])) {
+    return DotOperationType::kMatrixVector;
+  }
+  if (lhsShape.size() == 1 && rhsShape.size() == 2 &&
+      shapeMatches(lhsShape[0], rhsShape[0])) {
+    return DotOperationType::kVectorMatrix;
+  }
+  if (lhsShape.size() == 2 && rhsShape.size() == 2 &&
+      shapeMatches(lhsShape[1], rhsShape[0])) {
+    return DotOperationType::kMatrixMatrix;
+  }
+  return DotOperationType::kUnsupported;
+}
+
+SmallVector<Value, 2> getDotOpEmptyTensorDynSizes(OpBuilder& b, Location loc,
+                                                  Value lhs, Value rhs,
+                                                  DotOperationType type) {
+  SmallVector<Value, 2> dynShape;
+  switch (type) {
+    case DotOperationType::kMatrixMatrix: {
+      if (lhs.getType().cast<ShapedType>().isDynamicDim(0))
+        dynShape.push_back(b.create<tensor::DimOp>(loc, lhs, 0));
+      if (rhs.getType().cast<ShapedType>().isDynamicDim(1))
+        dynShape.push_back(b.create<tensor::DimOp>(loc, rhs, 1));
+      break;
+    }
+    case DotOperationType::kMatrixVector: {
+      if (lhs.getType().cast<ShapedType>().isDynamicDim(0))
+        dynShape.push_back(b.create<tensor::DimOp>(loc, lhs, 0));
+      break;
+    }
+    case DotOperationType::kVectorMatrix: {
+      if (rhs.getType().cast<ShapedType>().isDynamicDim(1))
+        dynShape.push_back(b.create<tensor::DimOp>(loc, rhs, 1));
+      break;
+    }
+    case DotOperationType::kVectorDot:
+    case DotOperationType::kUnsupported:
+      break;
+  }
+  return dynShape;
+}
+
+template <DotOperationType op_type, typename LinalgOp>
+class DotOpConversion : public OpConversionPattern<stablehlo::DotOp> {
+ public:
+  using OpConversionPattern<stablehlo::DotOp>::OpConversionPattern;
+  LogicalResult matchAndRewrite(
+      stablehlo::DotOp op, stablehlo::DotOp::Adaptor adaptor,
+      ConversionPatternRewriter& rewriter) const final {
+    if (!verifyHloOpBufferOrTensorSemantics(op)) {
+      return failure();
+    }
+    if (getDotOperationType(op) != op_type) return failure();
+
+    Location loc = op.getLoc();
+    // Convert unsigned to signed. This works because signed and unsigned
+    // integer matmul is the same operation in two's complement.
+    auto outputType =
+        typeConverter->convertType(op.getType()).cast<ShapedType>();
+    SmallVector<Value, 2> dynShape = getDotOpEmptyTensorDynSizes(
+        rewriter, loc, adaptor.getLhs(), adaptor.getRhs(), op_type);
+    auto emptyTensor =
+        sparse_tensor::getSparseTensorEncoding(outputType) == nullptr
+            ? getEmptyTensor(rewriter, loc, outputType, dynShape)
+            : getEmptySparseTensor(rewriter, loc, outputType, dynShape);
+    Value zeroTensor = fillTensorWithZeros(rewriter, loc, emptyTensor);
+    rewriter.replaceOpWithNewOp<LinalgOp>(
+        op, TypeRange{outputType},
+        ValueRange{adaptor.getLhs(), adaptor.getRhs()}, ValueRange{zeroTensor},
+        linalg::getPrunedAttributeList(op));
+    return success();
+  }
+};
+
+class DotGeneralBatchMatMulOpConversion
+    : public OpConversionPattern<stablehlo::DotGeneralOp> {
+ public:
+  using OpConversionPattern<stablehlo::DotGeneralOp>::OpConversionPattern;
+  LogicalResult matchAndRewrite(
+      stablehlo::DotGeneralOp op, OpAdaptor adaptor,
+      ConversionPatternRewriter& rewriter) const final {
+    if (!verifyHloOpBufferOrTensorSemantics(op)) {
+      return failure();
+    }
+    if (op.getType().cast<RankedTensorType>().getRank() != 3) {
+      return rewriter.notifyMatchFailure(op, "expected a batch matmul");
+    }
+
+    stablehlo::DotDimensionNumbersAttr dimNumbers = op.getDotDimensionNumbers();
+    auto lhsBatchingDims = dimNumbers.getLhsBatchingDimensions();
+    auto rhsBatchingDims = dimNumbers.getRhsBatchingDimensions();
+    auto lhsContractingDims = dimNumbers.getLhsContractingDimensions();
+    auto rhsContractingDims = dimNumbers.getRhsContractingDimensions();
+    if (lhsBatchingDims.size() != 1 || lhsBatchingDims[0] != 0) {
+      return rewriter.notifyMatchFailure(
+          op, "expected lhs batching dimensions exactly {0}");
+    }
+    if (rhsBatchingDims.size() != 1 || rhsBatchingDims[0] != 0) {
+      return rewriter.notifyMatchFailure(
+          op, "expected rhs batching dimensions exactly {0}");
+    }
+    if (lhsContractingDims.size() != 1 || lhsContractingDims[0] != 2) {
+      return rewriter.notifyMatchFailure(
+          op, "expected lhs contracting dimensions exactly {2}");
+    }
+    if (rhsContractingDims.size() != 1 || rhsContractingDims[0] != 1) {
+      return rewriter.notifyMatchFailure(
+          op, "expected rhs contracting dimensions exactly {1}");
+    }
+
+    Location loc = op.getLoc();
+    // Convert unsigned to signed. This works because signed and unsigned
+    // integer matmul is the same operation in two's complement.
+    auto outputType =
+        typeConverter->convertType(op.getType()).cast<ShapedType>();
+    auto emptyTensor =
+        getEmptyTensorFor(rewriter, loc, outputType, op, adaptor.getOperands());
+    Value zeroTensor = fillTensorWithZeros(rewriter, loc, emptyTensor);
+    Operation* linalgOp = rewriter.create<linalg::BatchMatmulOp>(
+        loc, /*resultTensorTypes=*/TypeRange{outputType},
+        /*inputs=*/ValueRange{adaptor.getLhs(), adaptor.getRhs()},
+        /*outputBuffers=*/ValueRange{zeroTensor},
+        linalg::getPrunedAttributeList(op));
+
+    rewriter.replaceOp(op, linalgOp->getResults());
+    return success();
+  }
+};
+
+class MapOpToGenericConverter : public OpConversionPattern<stablehlo::MapOp> {
+ public:
+  using OpConversionPattern::OpConversionPattern;
+  LogicalResult matchAndRewrite(
+      stablehlo::MapOp op, OpAdaptor adaptor,
+      ConversionPatternRewriter& rewriter) const final {
+    if (!verifyHloOpBufferOrTensorSemantics(op)) return failure();
+
+    auto resultType =
+        typeConverter->convertType(op.getType()).cast<ShapedType>();
+    assert(op.getDimensions().size() == resultType.getRank() &&
+           "Expected a pointwise map");
+
+    Location loc = op.getLoc();
+    Value output =
+        getEmptyTensorFor(rewriter, loc, resultType, op, adaptor.getOperands());
+    SmallVector<AffineMap> indexingMaps(
+        op.getNumOperands() + 1,
+        rewriter.getMultiDimIdentityMap(resultType.getRank()));
+
+    auto linalgOp = rewriter.create<linalg::GenericOp>(
+        loc, resultType, adaptor.getOperands(), output, indexingMaps,
+        getNParallelLoopsAttrs(resultType.getRank()),
+        /*bodyBuild=*/nullptr, linalg::getPrunedAttributeList(op));
+
+    // Convert the signature of the body. We scalarize the operands and add a
+    // scalar operand representing the output tensor.
+    Region& region = linalgOp.getRegion();
+    rewriter.inlineRegionBefore(op.getComputation(), region, region.end());
+    TypeConverter::SignatureConversion signatureConverter(op.getNumOperands() +
+                                                          1);
+
+    for (const auto& it : llvm::enumerate(op.getOperation()->getOperands())) {
+      signatureConverter.addInputs(
+          it.index(),
+          typeConverter->convertType(
+              it.value().getType().cast<ShapedType>().getElementType()));
+    }
+    signatureConverter.addInputs(resultType.getElementType());
+
+    rewriter.applySignatureConversion(&region, signatureConverter,
+                                      getTypeConverter());
+    rewriter.replaceOp(op, linalgOp.getResults());
+    return success();
+  }
+};
+
+class MapOpToMapConverter : public OpConversionPattern<stablehlo::MapOp> {
+ public:
+  using OpConversionPattern::OpConversionPattern;
+  LogicalResult matchAndRewrite(
+      stablehlo::MapOp op, OpAdaptor adaptor,
+      ConversionPatternRewriter& rewriter) const final {
+    if (!verifyHloOpBufferOrTensorSemantics(op)) return failure();
+
+    auto resultType =
+        typeConverter->convertType(op.getType()).cast<ShapedType>();
+    assert(op.getDimensions().size() == resultType.getRank() &&
+           "Expected a pointwise map");
+
+    Location loc = op.getLoc();
+    Value operand0 = adaptor.getOperands()[0];
+    SmallVector<Value> coercedOperands = {operand0};
+    for (Value operand : llvm::drop_begin(adaptor.getOperands(), 1)) {
+      coercedOperands.push_back(coerceTensorShape(
+          rewriter, loc, cast<TypedValue<ShapedType>>(operand),
+          operand0.getType()));
+    }
+    Value output = rewriter.create<tensor::EmptyOp>(
+        loc, tensor::getMixedSizes(rewriter, loc, operand0),
+        resultType.getElementType());
+
+    auto linalgOp = rewriter.create<linalg::MapOp>(
+        loc, coercedOperands, output,
+        /*bodyBuild=*/nullptr, linalg::getPrunedAttributeList(op));
+
+    // Convert the signature of the body. We scalarize the operands and add a
+    // scalar operand representing the output tensor.
+    Region& region = linalgOp.getRegion();
+    rewriter.inlineRegionBefore(op.getComputation(), region, region.end());
+    TypeConverter::SignatureConversion signatureConverter(op.getNumOperands());
+
+    for (const auto& it : llvm::enumerate(op.getOperation()->getOperands())) {
+      signatureConverter.addInputs(
+          it.index(),
+          typeConverter->convertType(
+              it.value().getType().cast<ShapedType>().getElementType()));
+    }
+
+    rewriter.applySignatureConversion(&region, signatureConverter,
+                                      getTypeConverter());
+    auto result = rewriter.createOrFold<tensor::CastOp>(loc, resultType,
+                                                        linalgOp.getResults());
+    rewriter.replaceOp(op, result);
+    return success();
+  }
+};
+
+SmallVector<Value, 8> getReduceOpEmptyTensorDynSizes(
+    OpBuilder& b, Location loc, Value arg, ShapedType resultType,
+    ArrayRef<int64_t> reductionDims) {
+  llvm::SmallSetVector<int, 4> s;
+  for (auto dim : reductionDims) s.insert(dim);
+
+  SmallVector<unsigned, 4> parallelDims;
+  SmallVector<Value, 8> dynShape;
+  int rank = arg.getType().cast<RankedTensorType>().getRank();
+  for (int i = 0, j = 0; i < rank; ++i) {
+    if (s.count(i)) continue;
+    if (!resultType.isDynamicDim(j++)) continue;
+    dynShape.push_back(b.create<tensor::DimOp>(loc, arg, i));
+  }
+
+  return dynShape;
+}
+
+class ReduceRegionReturnOpConversion
+    : public OpConversionPattern<stablehlo::ReturnOp> {
+ public:
+  using OpConversionPattern<stablehlo::ReturnOp>::OpConversionPattern;
+  LogicalResult matchAndRewrite(
+      stablehlo::ReturnOp op, OpAdaptor adaptor,
+      ConversionPatternRewriter& rewriter) const final {
+    if (!isInBodyOfLinalgOps(op)) {
+      return failure();
+    }
+    SmallVector<Value, 4> operands(adaptor.getOperands());
+    for (size_t i = 0; i < operands.size(); ++i) {
+      if (operands[i].getType().isa<ShapedType>()) {
+        auto loc = operands[i].getLoc();
+        operands[i] = rewriter.create<tensor::ExtractOp>(loc, operands[i]);
+      }
+    }
+    rewriter.replaceOpWithNewOp<linalg::YieldOp>(op, operands);
+    return success();
+  }
+};
+
+class ReduceOpToGenericConverter
+    : public OpConversionPattern<stablehlo::ReduceOp> {
+ public:
+  using OpConversionPattern<stablehlo::ReduceOp>::OpConversionPattern;
+  LogicalResult matchAndRewrite(
+      stablehlo::ReduceOp op, OpAdaptor adaptor,
+      ConversionPatternRewriter& rewriter) const final {
+    Location loc = op.getLoc();
+
+    int numOperands = static_cast<int>(adaptor.getInputs().size());
+
+    if (llvm::any_of(adaptor.getInputs(), [](Value v) {
+          return !v.getType().isa<RankedTensorType>();
+        })) {
+      return rewriter.notifyMatchFailure(op, "expects known-rank args");
+    }
+    auto srcRank =
+        adaptor.getInputs()[0].getType().cast<ShapedType>().getRank();
+
+    SmallVector<int64_t, 4> reductionDims = extract1DVector(op.getDimensions());
+
+    SmallVector<Type> resultTypes;
+    if (failed(typeConverter->convertTypes(op.getResultTypes(), resultTypes)))
+      return failure();
+
+    SmallVector<Value> outputs;
+    SmallVector<AffineMap, 3> indexingMaps;
+    for (auto [operand, initValue, resultType] :
+         llvm::zip(adaptor.getInputs(), adaptor.getInitValues(), resultTypes)) {
+      // Check if init_value is constant. If so, inline the value into the
+      // region.
+      initValue = rewriter.createOrFold<tensor::ExtractOp>(loc, initValue);
+
+      SmallVector<Value, 8> dynShape = getReduceOpEmptyTensorDynSizes(
+          rewriter, loc, operand, resultType, reductionDims);
+      auto emptyTensor = getEmptyTensor(rewriter, loc, resultType, dynShape);
+      Value filledTensor =
+          rewriter.create<linalg::FillOp>(loc, initValue, emptyTensor).result();
+      outputs.push_back(filledTensor);
+    }
+
+    // Prepare indexing maps for linalg generic op. The elements are for src
+    // and dst. Transpose `src` to make the reduction loops be the innermost,
+    // because it's easier to fully utilize processors.
+    indexingMaps.append(
+        numOperands, getTransposeMapForReduction(rewriter.getContext(),
+                                                 (int)srcRank, reductionDims));
+
+    // The indexing map of `dst` should drop the reduction loops. Since the
+    // reduction loops now are all in the innermost, drops
+    // `reduction_dims.size()` dimensions. We don't need an inverse
+    // permutation here because they are the same.
+    SmallVector<AffineExpr, 4> exprs;
+    for (int i = 0, e = srcRank - reductionDims.size(); i < e; ++i)
+      exprs.push_back(rewriter.getAffineDimExpr(i));
+    indexingMaps.append(numOperands,
+                        AffineMap::get(srcRank, /*symbolCount=*/0, exprs,
+                                       rewriter.getContext()));
+
+    auto linalgOp = rewriter.create<linalg::GenericOp>(
+        loc, /*resultTensorTypes=*/resultTypes, adaptor.getInputs(),
+        /*outputBuffers=*/ValueRange{outputs}, indexingMaps,
+        getParallelAndReductionIterators(srcRank, reductionDims.size()),
+        /*bodyBuild=*/nullptr, linalg::getPrunedAttributeList(op));
+
+    // Convert the signature of the body. The reduce op region apply function
+    // has a signature (lhs, rhs) -> output, all of the same tensor type t.
+    // This is converted to a function with the same signature but with
+    // element types. E.g., "(tensor<f32>, tensor<f32>) -> tensor<f32>" will
+    // be converted to "(f32, f32, f32)".
+    Region& region = linalgOp.getRegion();
+    rewriter.inlineRegionBefore(op.getBody(), region, region.end());
+    TypeConverter::SignatureConversion signatureConverter(numOperands * 2);
+
+    // The mhlo ReduceOp requires that the seed be used as a LHS operand inside
+    // the region, and the seed is encoded in linalg in the intial out value, so
+    // modify the signature of the block and the value mappings, so the output
+    // args will correlate with the original LHS and the inputs correlate with
+    // the original RHS.
+    for (const auto& [idx, val] : llvm::enumerate(op.getInputs())) {
+      signatureConverter.addInputs(
+          /*origInputNo=*/idx + numOperands,
+          // type for the new operand number 'idx'.
+          typeConverter->convertType(
+              val.getType().cast<ShapedType>().getElementType()));
+    }
+    for (const auto& [idx, val] : llvm::enumerate(op.getInitValues())) {
+      signatureConverter.addInputs(
+          /*origInputNo=*/idx,
+          // type for the new operand number 'idx' + 'numOperands'.
+          typeConverter->convertType(
+              val.getType().cast<ShapedType>().getElementType()));
+    }
+
+    rewriter.applySignatureConversion(&region, signatureConverter,
+                                      getTypeConverter());
+    rewriter.replaceOp(op, linalgOp.getResults());
+    return success();
+  }
+};
+
+struct ReduceOpToReduceConverter
+    : public OpConversionPattern<stablehlo::ReduceOp> {
+  using OpConversionPattern<stablehlo::ReduceOp>::OpConversionPattern;
+
+  LogicalResult matchAndRewrite(
+      stablehlo::ReduceOp op, OpAdaptor adaptor,
+      ConversionPatternRewriter& rewriter) const final {
+    auto reductionDims =
+        llvm::to_vector(op.getDimensions().getValues<int64_t>());
+    // mhlo.reduce doesn't specify the order of the reduction dimensions.
+    llvm::sort(reductionDims);
+
+    auto toRankedTensor = [](Value v) -> RankedTensorType {
+      return v.getType().dyn_cast<RankedTensorType>();
+    };
+
+    SmallVector<Value> outputs;
+    SmallVector<RankedTensorType> operandTypes, initTypes;
+    SmallVector<Type> resultTypes;
+    if (failed(typeConverter->convertTypes(op.getResultTypes(), resultTypes)))
+      return failure();
+
+    Location loc = op.getLoc();
+    for (auto [operand, initValue, resultType] :
+         llvm::zip(adaptor.getInputs(), adaptor.getInitValues(), resultTypes)) {
+      auto initType = toRankedTensor(initValue);
+      if (!initType)
+        return rewriter.notifyMatchFailure(op,
+                                           "expects known-rank init values");
+      initTypes.push_back(initType);
+      auto operandType = toRankedTensor(operand);
+      if (!operandType)
+        return rewriter.notifyMatchFailure(op, "expects known-rank operands");
+      operandTypes.push_back(operandType);
+      initValue = rewriter.createOrFold<tensor::ExtractOp>(loc, initValue);
+      auto tensorResultType = resultType.cast<RankedTensorType>();
+      // For linalg.reduce, the result type's dimensions must match the input's
+      // dimensions, whereas MHLO allows replacing static dimensions with
+      // dynamic ones.
+      SmallVector<int64_t> resultShape;
+      SmallVector<Value, 8> dynShape;
+      for (auto [index, dim] :
+           llvm::enumerate(operand.getType().cast<ShapedType>().getShape())) {
+        if (!llvm::is_contained(reductionDims, index)) {
+          resultShape.push_back(dim);
+          if (ShapedType::isDynamic(dim)) {
+            dynShape.push_back(
+                rewriter.create<tensor::DimOp>(loc, operand, index));
+          }
+        }
+      }
+
+      Value emptyTensor = rewriter.create<tensor::EmptyOp>(
+          loc, resultShape, tensorResultType.getElementType(), dynShape);
+      Value filledTensor =
+          rewriter.create<linalg::FillOp>(loc, initValue, emptyTensor).result();
+      outputs.push_back(filledTensor);
+    }
+
+    auto linalgOp = rewriter.create<linalg::ReduceOp>(
+        loc, adaptor.getInputs(), outputs, reductionDims,
+        /*bodyBuild=*/nullptr, linalg::getPrunedAttributeList(op));
+
+    Region& region = linalgOp.getRegion();
+    rewriter.inlineRegionBefore(op.getBody(), region, region.end());
+
+    // Convert the signature of the body. The reduce op 'computation' region
+    // apply function has a signature with tensor types, this is converted to a
+    // function with element types. E.g. the signature "(tensor<f32>,
+    // tensor<f32>) -> tensor<f32>" will be converted to "(f32, f32) -> f32".
+    // Also, we need to swap the operands of the function. The mhlo.reduce op
+    // expects the init values to be the first parameters of the apply function,
+    // while the linalg.reduction op expects the init values as the last
+    // parameters of the 'combiner' region apply function.
+    TypeConverter::SignatureConversion signatureConverter(
+        linalgOp.getNumDpsInputs() * 2);
+    assert(linalgOp.getNumDpsInputs() == linalgOp.getNumDpsInits());
+    for (const auto& [idx, val] : llvm::enumerate(operandTypes)) {
+      signatureConverter.addInputs(
+          /*origInputNo=*/idx + linalgOp.getNumDpsInputs(),
+          // type for new operand number 'idx'.
+          typeConverter->convertType(val.getElementType()));
+    }
+    for (const auto& [idx, val] : llvm::enumerate(initTypes)) {
+      signatureConverter.addInputs(
+          /*origInputNo=*/idx,
+          // type for new operand number 'idx' + linalgOp.getNumInputs()
+          typeConverter->convertType(val.getElementType()));
+    }
+    rewriter.applySignatureConversion(&region, signatureConverter,
+                                      getTypeConverter());
+
+    // Cast the result to the correct type.
+    SmallVector<Value> results;
+    for (auto [result, resultType] :
+         llvm::zip(linalgOp.getResults(), resultTypes)) {
+      results.push_back(
+          rewriter.createOrFold<tensor::CastOp>(loc, resultType, result));
+    }
+    rewriter.replaceOp(op, results);
+    return success();
+  }
+};
+
+/// Converts xla-hlo.select_and_scatter op to a sequence of linalg.generics ops.
+/// The current version computes the scattered index and populates the correct
+/// value for each tile. It does not currently handle overlapping tiles.
+struct SelectAndScatterNoOverlapConverter
+    : public OpConversionPattern<stablehlo::SelectAndScatterOp> {
+  using OpConversionPattern<stablehlo::SelectAndScatterOp>::OpConversionPattern;
+
+  LogicalResult matchAndRewrite(
+      stablehlo::SelectAndScatterOp op, OpAdaptor adaptor,
+      ConversionPatternRewriter& rewriter) const final {
+    Location loc = op.getLoc();
+    ImplicitLocOpBuilder b(loc, rewriter);
+    Value source = op.getSource();
+    Value operand = op.getOperand();
+    Value init = op.getInitValue();
+
+    auto sourceTy = source.getType().dyn_cast<RankedTensorType>();
+    auto operandTy = operand.getType().dyn_cast<RankedTensorType>();
+    auto initTy = init.getType().dyn_cast<RankedTensorType>();
+    auto resultTy = op.getResult().getType().dyn_cast<RankedTensorType>();
+    if (!sourceTy || !operandTy || !initTy || !resultTy)
+      return rewriter.notifyMatchFailure(op, "inputs/outputs must be ranked");
+
+    auto indexETy = b.getI32Type();
+    auto srcETy = operandTy.getElementType();
+    auto destETy = initTy.getElementType();
+
+    const int64_t rank = sourceTy.getRank();
+
+    llvm::SmallVector<int64_t> pad(rank * 2, 0);
+    if (op.getPadding().has_value())
+      pad = llvm::to_vector(op.getPaddingAttr().getValues<int64_t>());
+
+    // TODO(suderman): Add support for padding.
+    if (llvm::any_of(pad, [](int64_t p) { return p != 0; }))
+      return rewriter.notifyMatchFailure(op, "non-zero padding values found.");
+
+    if (!op.getWindowStrides().has_value())
+      return rewriter.notifyMatchFailure(op, "no window strides found");
+
+    if (!op.getWindowDimensions().has_value())
+      return rewriter.notifyMatchFailure(op, "no window dimensions found");
+
+    auto strides =
+        llvm::to_vector(op.getWindowStridesAttr().getValues<int64_t>());
+    auto window =
+        llvm::to_vector(op.getWindowDimensionsAttr().getValues<int64_t>());
+
+    if (static_cast<int64_t>(strides.size()) != operandTy.getRank() ||
+        static_cast<int64_t>(window.size()) != operandTy.getRank())
+      return rewriter.notifyMatchFailure(
+          op, "stride/window length should equal operand rank");
+
+    // The current version cannot handle overlapped regions.
+    for (int i = 0, s = strides.size(); i < s; ++i) {
+      if (strides[i] < window[i])
+        return rewriter.notifyMatchFailure(
+            op, "overlapping windows are not supported");
+    }
+
+    // If the window only contains a single element, this lowering will be
+    // problematic. Ultimately we should handle this with a canonicalizer.
+    if (llvm::all_of(window, [](auto sz) { return sz == 1; })) {
+      return rewriter.notifyMatchFailure(op,
+                                         "unary window size is not supported");
+    }
+
+    // The first linalg.generic operation computes the relevant index over
+    // window for the defined mhlo.select_and_scatter. This involves
+    // iterating over the window of the operand a computing the index.
+    // Rather than storing N indices we compute the row major identifier
+    // in the window, to specify which location should be scattered to.
+
+    // Output specifies the `rank` parallel iterators that correspond to
+    // output values.
+    SmallVector<AffineExpr, 4> outputExprs;
+    for (int i = 0, s = rank; i < s; ++i)
+      outputExprs.push_back(b.getAffineDimExpr(i));
+
+    // For the output we need to define the reduction across the window
+    // width and height. This includes applying striding behavior and
+    // adding the additional reduction iterators. We skip length-1 dimensions
+    // as the reduction is degenerate.
+    SmallVector<int64_t> filteredWindows, filteredStrides;
+    SmallVector<AffineExpr, 4> sourceExprs(outputExprs);
+    SmallVector<AffineExpr, 4> windowExprs;
+    for (int i = 0, s = rank; i < s; ++i) {
+      sourceExprs[i] = sourceExprs[i] * strides[i];
+      if (strides[i] != 1) {
+        auto expr = b.getAffineDimExpr(windowExprs.size() + sourceExprs.size());
+        sourceExprs[i] = sourceExprs[i] + expr;
+        windowExprs.push_back(expr);
+        filteredWindows.push_back(window[i]);
+        filteredStrides.push_back(strides[i]);
+      }
+    }
+
+    // Determine the total number of AffineExprs and construct the IndexingMaps
+    // for the windowed reduction operation.
+    const int64_t reduceExprCount = windowExprs.size() + sourceExprs.size();
+    SmallVector<AffineMap, 2> reduceIndexingMaps;
+    reduceIndexingMaps.push_back(AffineMap::get(reduceExprCount,
+                                                /*symbolCount=*/0, sourceExprs,
+                                                rewriter.getContext()));
+    reduceIndexingMaps.push_back(AffineMap::get(reduceExprCount,
+                                                /*symbolCount=*/0, windowExprs,
+                                                rewriter.getContext()));
+    auto reduceOutMap =
+        AffineMap::get(reduceExprCount,
+                       /*symbolCount=*/0, outputExprs, rewriter.getContext());
+    reduceIndexingMaps.push_back(reduceOutMap);
+    reduceIndexingMaps.push_back(reduceOutMap);
+
+    // Output sizes should match the dimensions of the `source` tensor, even if
+    // dynamic.
+    SmallVector<Value, 4> reduceDynSizes;
+    for (int i = 0, s = rank; i < s; ++i)
+      if (sourceTy.isDynamicDim(i))
+        reduceDynSizes.push_back(b.create<tensor::DimOp>(source, i));
+
+    Value reduceValueEmpty =
+        b.create<tensor::EmptyOp>(sourceTy.getShape(), destETy, reduceDynSizes);
+    Value reduceIndexEmpty = b.create<tensor::EmptyOp>(
+        sourceTy.getShape(), indexETy, reduceDynSizes);
+
+    // We initialize indices to -1 which indicates no matching destination.
+    Value negativeOne = b.create<arith::ConstantOp>(b.getI32IntegerAttr(-1));
+    reduceIndexEmpty =
+        b.create<linalg::FillOp>(negativeOne, reduceIndexEmpty).getResult(0);
+
+    // We only care to match the reduction dimensions.
+    Value windowEmpty = b.create<tensor::EmptyOp>(filteredWindows, srcETy);
+
+    auto reduceGeneric = b.create<linalg::GenericOp>(
+        /*resultTensors=*/ArrayRef<Type>{reduceValueEmpty.getType(),
+                                         reduceIndexEmpty.getType()},
+        /*inputs=*/ValueRange{operand, windowEmpty},
+        /*outputs=*/ValueRange{reduceValueEmpty, reduceIndexEmpty},
+        reduceIndexingMaps,
+        getParallelAndReductionIterators(reduceExprCount, windowExprs.size()),
+        /*bodyBuild=*/nullptr, linalg::getPrunedAttributeList(op));
+
+    // First we clone in the selection block.
+    auto& reduceRegion = reduceGeneric.getRegion();
+    rewriter.setInsertionPoint(reduceGeneric);
+    rewriter.cloneRegionBefore(op.getSelect(), reduceRegion,
+                               reduceRegion.end());
+
+    // This includes convert `mhlo` scalar-tensor regions to `linalg` scalars.
+    TypeConverter::SignatureConversion reduceSignConverter(4);
+    reduceSignConverter.addInputs(0, srcETy);
+    reduceSignConverter.addInputs(srcETy);
+    reduceSignConverter.addInputs(1, destETy);
+    reduceSignConverter.addInputs(indexETy);
+    rewriter.applySignatureConversion(&reduceRegion, reduceSignConverter,
+                                      getTypeConverter());
+
+    // Grab the terminator and use the turned value to now select the
+    // correct index and value.
+    auto& reduceBlock = reduceRegion.front();
+    auto* reduceTerminator = reduceBlock.getTerminator();
+    Value selectPred = reduceTerminator->getOperand(0);
+    Value selectInVal = reduceBlock.getArgument(0);
+    Value selectOutVal = reduceBlock.getArgument(2);
+    Value selectOutIdx = reduceBlock.getArgument(3);
+
+    b.setInsertionPoint(reduceTerminator);
+
+    // The predicate operates on scalar-tensors, so we need to extract the
+    // value for `linalg` operations. Tensor-ops are cleaned up by other
+    // rewriters.
+    selectPred = b.create<tensor::ExtractOp>(rewriter.getI1Type(), selectPred,
+                                             ValueRange{});
+
+    // We select if either the selection function returns `true` or the
+    // current reduction index is `-1`, e.g. no index has been selected yet.
+    Value selectNegOne = b.create<arith::CmpIOp>(arith::CmpIPredicate::eq,
+                                                 selectOutIdx, negativeOne);
+    selectPred = b.create<arith::OrIOp>(selectPred, selectNegOne);
+
+    // We compute a unique idx for each element in the window.
+    Value computedIdx = b.create<linalg::IndexOp>(rank);
+    for (int i = 1, s = filteredStrides.size(); i < s; ++i) {
+      Value width = b.create<arith::ConstantIndexOp>(filteredStrides[i]);
+      Value idx = b.create<linalg::IndexOp>(rank + i);
+      computedIdx = b.create<arith::MulIOp>(width, computedIdx);
+      computedIdx = b.create<arith::AddIOp>(computedIdx, idx);
+    }
+    computedIdx = b.create<arith::IndexCastOp>(indexETy, computedIdx);
+
+    // Using the selection predicate track the value and selected
+    // identifier for the future scattering.
+    Value selectedIdx =
+        b.create<arith::SelectOp>(selectPred, computedIdx, selectOutIdx);
+    Value selectedValue =
+        b.create<arith::SelectOp>(selectPred, selectInVal, selectOutVal);
+    b.create<linalg::YieldOp>(ValueRange{selectedValue, selectedIdx});
+
+    // Original terminator is an mhlo.return we no longer need.
+    rewriter.eraseOp(reduceTerminator);
+    b.setInsertionPoint(op);
+
+    Value reduceIndex = reduceGeneric.getResult(1);
+    ShapedType reduceIndexTy = reduceIndex.getType().cast<ShapedType>();
+
+    // For the second generic we restricted to only cases where there are
+    // no window overlaps. This guarantees that each source value is scattered
+    // within its own unique window. We can broadcast to this window size and
+    // populate only the relative location.
+    llvm::SmallVector<int64_t> broadcastShape;
+    llvm::SmallVector<Value> broadcastDynDims;
+    llvm::SmallVector<AffineExpr> broadcastExprs;
+    for (int i = 0, s = reduceIndexTy.getRank(); i < s; ++i) {
+      int64_t broadcast = strides[i];
+      if (sourceTy.isDynamicDim(i))
+        broadcastDynDims.push_back(b.create<tensor::DimOp>(source, i));
+
+      broadcastExprs.push_back(b.getAffineDimExpr(broadcastShape.size()));
+      broadcastShape.push_back(sourceTy.getDimSize(i));
+      if (broadcast > 1) {
+        broadcastShape.push_back(broadcast);
+      }
+    }
+
+    // We broadcast the values of our input tensors across the stride-tiling
+    // size.
+    Value scatterEmpty = b.create<tensor::EmptyOp>(
+        broadcastShape, resultTy.getElementType(), broadcastDynDims);
+    Value initScalar = b.create<tensor::ExtractOp>(initTy.getElementType(),
+                                                   init, ValueRange{});
+    Value scatterFill =
+        b.create<linalg::FillOp>(initScalar, scatterEmpty).getResult(0);
+
+    // Both the indices and values are broadcasted using the same indexing map.
+    // Output fully parallel.
+    auto scatterInputMap =
+        AffineMap::get(broadcastShape.size(), /*symbolCount=*/0, broadcastExprs,
+                       b.getContext());
+    SmallVector<AffineMap> scatterIndexingMaps;
+    scatterIndexingMaps.push_back(scatterInputMap);
+    scatterIndexingMaps.push_back(scatterInputMap);
+    scatterIndexingMaps.push_back(
+        b.getMultiDimIdentityMap(broadcastShape.size()));
+
+    auto scatterGeneric = b.create<linalg::GenericOp>(
+        /*resultTensors=*/ArrayRef<Type>{scatterFill.getType()},
+        /*inputs=*/ValueRange{reduceIndex, source},
+        /*outputs=*/ValueRange{scatterFill}, scatterIndexingMaps,
+        getNParallelLoopsAttrs(broadcastShape.size()),
+        /*bodyBuild=*/nullptr, linalg::getPrunedAttributeList(op));
+
+    // Clone the scattering combination logic and perform the tensor-to-scalar
+    // conversion.
+    auto& scatterRegion = scatterGeneric.getRegion();
+    b.setInsertionPoint(scatterGeneric);
+    rewriter.cloneRegionBefore(op.getScatter(), scatterRegion,
+                               scatterRegion.end());
+
+    TypeConverter::SignatureConversion scatterSignConverter(4);
+    scatterSignConverter.addInputs(indexETy);
+    scatterSignConverter.addInputs(0, sourceTy.getElementType());
+    scatterSignConverter.addInputs(1, sourceTy.getElementType());
+    rewriter.applySignatureConversion(&scatterRegion, scatterSignConverter,
+                                      getTypeConverter());
+
+    auto& scatterBlock = scatterRegion.front();
+    auto scatterTerminator = scatterBlock.getTerminator();
+    b.setInsertionPoint(scatterTerminator);
+
+    Value scatterInputIdx = scatterBlock.getArgument(0);
+    Value scatterOutputVal = scatterBlock.getArgument(2);
+    Value scatterUpdate = b.create<tensor::ExtractOp>(
+        sourceTy.getElementType(), scatterTerminator->getOperand(0),
+        ValueRange{});
+
+    // Compute the index of the tiled region to determine if it was selected.
+    Value id = b.create<arith::ConstantIndexOp>(0);
+    int64_t dim = 0;
+    for (int i = 0, s = strides.size(); i < s; ++i) {
+      if (strides[i] > 1) {
+        Value idx = b.create<linalg::IndexOp>(++dim);
+        Value tileSz = b.create<arith::ConstantIndexOp>(strides[i]);
+        id = b.create<arith::MulIOp>(id, tileSz);
+        id = b.create<arith::AddIOp>(id, idx);
+      }
+      ++dim;
+    }
+
+    // Check whether the computed id matches the to-scatter id, then select and
+    // yield.
+    id = b.create<arith::IndexCastOp>(indexETy, id);
+    auto scatterPred = b.create<arith::CmpIOp>(
+        b.getI1Type(), arith::CmpIPredicate::eq, id, scatterInputIdx);
+    scatterUpdate =
+        b.create<arith::SelectOp>(scatterPred, scatterUpdate, scatterOutputVal);
+
+    b.create<linalg::YieldOp>(scatterUpdate);
+    rewriter.eraseOp(scatterTerminator);
+    b.setInsertionPoint(op);
+
+    // We now need to collapse the tiles back into their
+    // source dimensions. We collapse any of the broadcast regions together.
+    int64_t collapseDim = 0;
+    SmallVector<ReassociationIndices> reassociationMap;
+    for (int i = 0, s = window.size(); i < s; ++i) {
+      SmallVector<int64_t, 2> dims = {collapseDim};
+      if (strides[i] > 1) dims.push_back(collapseDim + 1);
+
+      reassociationMap.push_back(ReassociationIndices(dims));
+      collapseDim += dims.size();
+    }
+
+    Value collapse = b.create<tensor::CollapseShapeOp>(
+        scatterGeneric.getResult(0), reassociationMap);
+    auto collapseTy = collapse.getType().cast<ShapedType>();
+
+    // After collapsing it it possible that the target may need to be padded.
+    auto zero = b.createOrFold<arith::ConstantIndexOp>(0);
+    SmallVector<int64_t> padShape;
+    SmallVector<OpFoldResult> padLow, padHigh;
+    padLow.resize(operandTy.getRank(), zero);
+
+    for (int i = 0, s = rank; i < s; ++i) {
+      int64_t size = std::max(resultTy.getDimSize(i), collapseTy.getDimSize(i));
+      if (operandTy.isDynamicDim(i) || collapseTy.isDynamicDim(i))
+        size = ShapedType::kDynamic;
+      padShape.push_back(size);
+
+      Value in = b.create<tensor::DimOp>(collapse, i);
+      Value out = b.create<tensor::DimOp>(operand, i);
+      Value diff = b.create<arith::SubIOp>(out, in);
+      Value pad = b.createOrFold<arith::MaxSIOp>(diff, zero);
+      padHigh.push_back(pad);
+    }
+
+    Value padded = b.create<tensor::PadOp>(collapseTy.clone(padShape), collapse,
+                                           padLow, padHigh, initScalar);
+
+    // The result may exceed the target size, slice if necessary.
+    SmallVector<OpFoldResult> sliceSizes;
+    SmallVector<OpFoldResult> sliceOffsets(operandTy.getRank(),
+                                           b.getIndexAttr(0));
+    SmallVector<OpFoldResult> sliceStrides(operandTy.getRank(),
+                                           b.getIndexAttr(1));
+    for (int i = 0, s = operandTy.getRank(); i < s; ++i) {
+      OpFoldResult dim = b.getIndexAttr(operandTy.getDimSize(i));
+      if (operandTy.isDynamicDim(i))
+        dim = b.createOrFold<tensor::DimOp>(operand, i);
+      sliceSizes.push_back(dim);
+    }
+
+    rewriter.setInsertionPoint(op);
+    rewriter.replaceOpWithNewOp<tensor::ExtractSliceOp>(
+        op, padded, sliceOffsets, sliceSizes, sliceStrides);
+
+    return success();
+  }
+};
+
+// Decomposes a pad with negative edge padding into a pad without negative edge
+// padding and a tensor.extract_slice.
+struct PadOpNegativePaddingConversion
+    : public OpConversionPattern<stablehlo::PadOp> {
+  using OpConversionPattern::OpConversionPattern;
+
+  LogicalResult matchAndRewrite(
+      stablehlo::PadOp op, OpAdaptor adaptor,
+      ConversionPatternRewriter& rewriter) const override {
+    SmallVector<int64_t, 4> padLow;
+    SmallVector<int64_t, 4> padHigh;
+    SmallVector<OpFoldResult, 4> sliceStarts;
+
+    bool hasNegativePadding = false;
+    for (int64_t low : op.getEdgePaddingLow().getValues<int64_t>()) {
+      if (low >= 0) {
+        padLow.push_back(low);
+        sliceStarts.push_back(rewriter.getIndexAttr(0));
+      } else {
+        padLow.push_back(0);
+        sliceStarts.push_back(rewriter.getIndexAttr(-low));
+        hasNegativePadding = true;
+      }
+    }
+
+    for (int64_t high : op.getEdgePaddingHigh().getValues<int64_t>()) {
+      if (high >= 0) {
+        padHigh.push_back(high);
+      } else {
+        padHigh.push_back(-high);
+        hasNegativePadding = true;
+      }
+    }
+
+    // If there's no negative edge padding we're done.
+    if (!hasNegativePadding) return failure();
+
+    // Create a new pad op with the positive values.
+    Value pad = rewriter.create<stablehlo::PadOp>(
+        op.getLoc(), adaptor.getOperand(), adaptor.getPaddingValue(),
+        rewriter.getI64TensorAttr(padLow), rewriter.getI64TensorAttr(padHigh),
+        op.getInteriorPadding());
+
+    // Then slice according to the negative edge padding. Static shapes only for
+    // now.
+    if (!op.getType().hasStaticShape()) return failure();
+    SmallVector<OpFoldResult, 4> sizes(llvm::map_range(
+        op.getType().getShape(),
+        [&](int64_t dim) { return rewriter.getIndexAttr(dim); }));
+    SmallVector<OpFoldResult, 4> strides(sliceStarts.size(),
+                                         rewriter.getIndexAttr(1));
+    rewriter.replaceOpWithNewOp<tensor::ExtractSliceOp>(op, pad, sliceStarts,
+                                                        sizes, strides);
+    return success();
+  }
+};
+
+/// Converts mhlo.pad operation to tensor.pad or tensor.insert_slice.
+struct PadOpConversion : public OpConversionPattern<stablehlo::PadOp> {
+  using OpConversionPattern::OpConversionPattern;
+
+  LogicalResult matchAndRewrite(
+      stablehlo::PadOp op, OpAdaptor adaptor,
+      ConversionPatternRewriter& rewriter) const override {
+    auto loc = op.getLoc();
+    auto resultType = typeConverter->convertType(op.getResult().getType());
+
+    // Negative edge padding is decomposed separately.
+    auto isNegative = [](const APInt& intVal) { return intVal.isNegative(); };
+    if (llvm::any_of(op.getEdgePaddingLow().getValues<APInt>(), isNegative) ||
+        llvm::any_of(op.getEdgePaddingHigh().getValues<APInt>(), isNegative))
+      return failure();
+
+    Value paddingVal = rewriter.createOrFold<tensor::ExtractOp>(
+        loc, adaptor.getPaddingValue());
+
+    SmallVector<OpFoldResult, 4> low(
+        op.getEdgePaddingLow().getValues<IntegerAttr>());
+
+    // If there is no interior padding lower to tensor.pad directly.
+    if (llvm::all_of(op.getInteriorPadding().getValues<APInt>(),
+                     [](const APInt& intVal) { return intVal.isZero(); })) {
+      SmallVector<OpFoldResult, 4> high(
+          op.getEdgePaddingHigh().getValues<IntegerAttr>());
+      auto padTensorOp = rewriter.create<tensor::PadOp>(
+          loc, resultType, adaptor.getOperand(), low, high, paddingVal);
+      rewriter.replaceOp(op, padTensorOp.getResult());
+      return success();
+    }
+
+    // We have interior padding, which can be lowered to tensor.insert_slice.
+    // Start by filling a result-sized tensor with the pad value.
+    auto emptyTensor =
+        getEmptyTensorFor(rewriter, loc, resultType, op, adaptor.getOperands());
+    auto fill =
+        rewriter.create<linalg::FillOp>(loc, paddingVal, emptyTensor).result();
+
+    // Get sizes of the original operand.
+    auto operandType = adaptor.getOperand().getType().cast<ShapedType>();
+    auto sizes = llvm::to_vector<4>(llvm::map_range(
+        llvm::seq<int64_t>(0, operandType.getRank()),
+        [&](int64_t dim) -> OpFoldResult {
+          if (!operandType.isDynamicDim(dim))
+            return rewriter.getIndexAttr(operandType.getDimSize(dim));
+          return rewriter.create<tensor::DimOp>(loc, adaptor.getOperand(), dim)
+              .getResult();
+        }));
+    // Map interior padding to strides.
+    auto strides = llvm::to_vector<4>(
+        llvm::map_range(op.getInteriorPadding().getValues<IntegerAttr>(),
+                        [&](IntegerAttr stride) -> OpFoldResult {
+                          return rewriter.getIntegerAttr(stride.getType(),
+                                                         stride.getValue() + 1);
+                        }));
+
+    rewriter.replaceOpWithNewOp<tensor::InsertSliceOp>(
+        op, adaptor.getOperand(), fill, low, sizes, strides);
+    return success();
+  }
+};
+
+/// Converts xla-hlo.torch_index_select op to a linalg.generic op.
+struct TorchIndexSelectOpConversion
+    : public OpConversionPattern<stablehlo::TorchIndexSelectOp> {
+  using OpConversionPattern::OpConversionPattern;
+
+  LogicalResult matchAndRewrite(
+      stablehlo::TorchIndexSelectOp op, OpAdaptor adaptor,
+      ConversionPatternRewriter& rewriter) const final {
+    int axis = static_cast<int>(op.getDim());
+    int batch = static_cast<int>(op.getBatchDims());
+    auto indexShapedType = adaptor.getIndex().getType().cast<ShapedType>();
+    int numIndices = static_cast<int>(indexShapedType.getRank());
+    auto operandShapedType = adaptor.getOperand().getType().cast<ShapedType>();
+    if (axis < 0) axis += static_cast<int>(operandShapedType.getRank());
+    if (batch < 0) batch += numIndices;
+
+    Location loc = op.getLoc();
+    auto resultType = this->typeConverter->convertType(op.getResult().getType())
+                          .cast<ShapedType>();
+    int rank = static_cast<int>(resultType.getRank());
+
+    // The output shape is
+    //   `params[:axis] + indices[batch_dims:] + params[axis + 1:]`
+    SmallVector<Value, 4> dynSizes;
+    for (int i = 0; i < rank; ++i) {
+      if (!resultType.isDynamicDim(i)) continue;
+      if (i < axis) {
+        dynSizes.push_back(
+            rewriter.create<tensor::DimOp>(loc, adaptor.getOperand(), i));
+      } else if (i < (axis + numIndices - batch)) {
+        int idx = i - axis + batch;
+        dynSizes.push_back(
+            rewriter.create<tensor::DimOp>(loc, adaptor.getIndex(), idx));
+      } else {
+        int idx = i - (axis + numIndices - batch) + axis + 1;
+        dynSizes.push_back(
+            rewriter.create<tensor::DimOp>(loc, adaptor.getOperand(), idx));
+      }
+    }
+
+    // Generate dummy tensor to preserve slice shape information.
+    SmallVector<int64_t> sliceShape;
+    SmallVector<Value, 4> dynSliceSizes;
+    SmallVector<AffineExpr, 4> sliceExprs;
+    auto resultShape = resultType.getShape();
+    for (int i = 0; i < axis; ++i) {
+      sliceExprs.push_back(rewriter.getAffineDimExpr(i));
+      sliceShape.push_back(resultShape[i]);
+      if (!resultType.isDynamicDim(i)) continue;
+      dynSliceSizes.push_back(
+          rewriter.create<tensor::DimOp>(loc, adaptor.getOperand(), i));
+    }
+    for (int i = axis + numIndices - batch; i < rank; ++i) {
+      sliceExprs.push_back(rewriter.getAffineDimExpr(i));
+      sliceShape.push_back(resultShape[i]);
+      if (!resultType.isDynamicDim(i)) continue;
+      int idx = i - (axis + numIndices - batch) + axis + 1;
+      dynSliceSizes.push_back(
+          rewriter.create<tensor::DimOp>(loc, adaptor.getOperand(), idx));
+    }
+
+    // Setup AffineMap for operand tensor.
+    SmallVector<AffineExpr, 4> exprs;
+    for (int i = 0; i < batch; ++i) {
+      exprs.push_back(rewriter.getAffineDimExpr(i));
+    }
+    for (int i = 0, e = numIndices - batch; i < e; ++i) {
+      exprs.push_back(rewriter.getAffineDimExpr(axis + i));
+    }
+
+    SmallVector<AffineMap, 2> indexingMaps;
+    indexingMaps.emplace_back(
+        AffineMap::get(rank, /*symbolCount=*/0, exprs, rewriter.getContext()));
+    indexingMaps.emplace_back(AffineMap::get(
+        rank, /*symbolCount=*/0, sliceExprs, rewriter.getContext()));
+    indexingMaps.emplace_back(rewriter.getMultiDimIdentityMap(rank));
+
+    Value sliceOp = rewriter.create<tensor::EmptyOp>(
+        loc, sliceShape, resultType.getElementType(), dynSliceSizes);
+
+    Value emptyOp = rewriter.create<tensor::EmptyOp>(
+        loc, resultType.getShape(), resultType.getElementType(), dynSizes);
+    auto linalgOp = rewriter.create<linalg::GenericOp>(
+        loc, /*resultTensors=*/ArrayRef<Type>{resultType},
+        /*inputs=*/ValueRange{adaptor.getIndex(), sliceOp},
+        /*outputs=*/emptyOp, indexingMaps, getNParallelLoopsAttrs(rank),
+        /*bodyBuild=*/nullptr, linalg::getPrunedAttributeList(op));
+
+    SmallVector<Type, 4> bodyArgTypes;
+    SmallVector<Value, 2> linalgOpArgs = {adaptor.getIndex(), sliceOp};
+    // Add a block to the region.
+    auto* region = &linalgOp.getRegion();
+    auto* block = rewriter.createBlock(region, region->end());
+    for (auto blockArgs : linalgOpArgs) {
+      bodyArgTypes.push_back(
+          blockArgs.getType().cast<ShapedType>().getElementType());
+    }
+    block->addArguments(bodyArgTypes,
+                        SmallVector<Location>(bodyArgTypes.size(), loc));
+    block->addArguments(resultType.getElementType(), loc);
+    OpBuilder::InsertionGuard guard(rewriter);
+    rewriter.setInsertionPointToEnd(block);
+
+    Value castedValue = rewriter.create<arith::IndexCastOp>(
+        loc, rewriter.getIndexType(), block->getArgument(0));
+
+    SmallVector<Value, 4> indices;
+    for (int i = 0; i < axis; ++i) {
+      indices.push_back(rewriter.create<linalg::IndexOp>(loc, i));
+    }
+    indices.push_back(castedValue);
+    for (int i = axis + numIndices - batch; i < rank; ++i) {
+      indices.push_back(rewriter.create<linalg::IndexOp>(loc, i));
+    }
+    Value res =
+        rewriter.create<tensor::ExtractOp>(loc, adaptor.getOperand(), indices);
+    rewriter.create<linalg::YieldOp>(loc, res);
+
+    rewriter.replaceOp(op, linalgOp.getResults());
+    return success();
+  }
+};
+
+class DotGeneralOpConversion
+    : public OpConversionPattern<stablehlo::DotGeneralOp> {
+ public:
+  using OpConversionPattern::OpConversionPattern;
+  LogicalResult matchAndRewrite(
+      stablehlo::DotGeneralOp op, OpAdaptor adaptor,
+      ConversionPatternRewriter& rewriter) const final {
+    if (!verifyHloOpBufferOrTensorSemantics(op)) {
+      return failure();
+    }
+
+    // Get various dimension iterator information
+    stablehlo::DotDimensionNumbersAttr dimNumbers = op.getDotDimensionNumbers();
+    auto lhsBatchingDims = dimNumbers.getLhsBatchingDimensions();
+    auto rhsBatchingDims = dimNumbers.getRhsBatchingDimensions();
+    auto lhsContractingDims = dimNumbers.getLhsContractingDimensions();
+    auto rhsContractingDims = dimNumbers.getRhsContractingDimensions();
+
+    // Get shape information and initialize output
+    assert(lhsContractingDims.size() == rhsContractingDims.size() &&
+           "number of contracting dims must be equal");
+    auto numContracting = lhsContractingDims.size();
+    // Convert unsigned to signed. This works because signed and unsigned
+    // integer matmul is the same operation in two's complement.
+    auto outputType =
+        typeConverter->convertType(op.getType()).cast<ShapedType>();
+    auto targetRank = outputType.getRank();
+    auto totalLoopCount = numContracting + targetRank;
+
+    auto lhsRank = adaptor.getLhs().getType().cast<ShapedType>().getRank();
+    auto lhsExtraDims =
+        lhsRank - lhsBatchingDims.size() - lhsContractingDims.size();
+    auto rhsRank = adaptor.getRhs().getType().cast<ShapedType>().getRank();
+
+    Location loc = op.getLoc();
+    auto emptyTensor =
+        getEmptyTensorFor(rewriter, loc, outputType, op, adaptor.getOperands());
+    Value zeroTensor = fillTensorWithZeros(rewriter, loc, emptyTensor);
+    SmallVector<AffineMap, 3> indexingMaps;
+
+    auto getMap = [&](int64_t rank, ArrayRef<int64_t> batchingDims,
+                      ArrayRef<int64_t> contractingDims, size_t extraDims) {
+      llvm::SmallVector<AffineExpr> indices(rank);
+      for (const auto& i : llvm::enumerate(batchingDims)) {
+        indices[i.value()] = rewriter.getAffineDimExpr(i.index());
+      }
+      for (const auto& i : llvm::enumerate(contractingDims)) {
+        indices[i.value()] = rewriter.getAffineDimExpr(i.index() + targetRank);
+      }
+      for (int i = 0; i < rank; ++i) {
+        if (!indices[i]) {
+          indices[i] = rewriter.getAffineDimExpr(extraDims++);
+        }
+      }
+      indexingMaps.push_back(AffineMap::get(/*dimCount=*/totalLoopCount,
+                                            /*symbolCount=*/0, indices,
+                                            op->getContext()));
+    };
+    getMap(lhsRank, lhsBatchingDims, lhsContractingDims,
+           lhsBatchingDims.size());
+    getMap(rhsRank, rhsBatchingDims, rhsContractingDims,
+           rhsBatchingDims.size() + lhsExtraDims);
+
+    {
+      SmallVector<AffineExpr, 4> dimExprs;
+      dimExprs.reserve(targetRank);
+      for (unsigned i = 0; i < targetRank; ++i)
+        dimExprs.push_back(rewriter.getAffineDimExpr(i));
+      indexingMaps.push_back(AffineMap::get(/*dimCount=*/totalLoopCount,
+                                            /*symbolCount=*/0, dimExprs,
+                                            op.getContext()));
+    }
+
+    Operation* linalgOp = rewriter.create<linalg::GenericOp>(
+        loc, /*resultTensorTypes=*/TypeRange{outputType},
+        /*inputs=*/ValueRange{adaptor.getLhs(), adaptor.getRhs()},
+        /*outputBuffers=*/ValueRange{zeroTensor}, indexingMaps,
+        getParallelAndReductionIterators(
+            /*nLoops=*/totalLoopCount,
+            /*nReduction=*/numContracting),
+        [](OpBuilder& b, Location loc, ValueRange) {
+          ImplicitLocOpBuilder builder(loc, b);
+          linalg::MatmulOp::regionBuilder(builder, *b.getInsertionBlock(), {});
+        },
+        linalg::getPrunedAttributeList(op));
+
+    rewriter.replaceOp(op, linalgOp->getResults());
+    return success();
+  }
+};
+
+/// Converts a HLO operation to a linalg.map op that contains the corresponding
+/// scalar operations.
+template <typename OpTy>
+class PointwiseToLinalgMapConverter : public OpConversionPattern<OpTy> {
+ public:
+  using OpConversionPattern<OpTy>::OpConversionPattern;
+
+  LogicalResult matchAndRewrite(
+      OpTy op, typename OpTy::Adaptor adaptor,
+      ConversionPatternRewriter& rewriter) const final {
+    auto loc = op.getLoc();
+    int64_t maxRank = getMaxRank(adaptor);
+
+    // Apply only if all operands are scalar or have the same rank. Some ops,
+    // like `mhlo.select`, support implicit broadcasting of scalars.
+    if (!llvm::all_of(adaptor.getOperands(), [&](Value v) {
+          int64_t r = getRank(v);
+          return r == 0 || r == maxRank;
+        })) {
+      return rewriter.notifyMatchFailure(
+          op, "Operands must be of same rank or scalar.");
+    }
+
+    // Find result type, if on tensors.
+    std::optional<ShapedType> resultTy;
+    resultTy = this->typeConverter->convertType(op->getResultTypes().front())
+                   .template dyn_cast<ShapedType>();
+
+    // Check result type compatibility.
+    if (!resultTy || !resultTy->hasRank() || resultTy->getRank() != maxRank ||
+        !(resultTy->getElementType().isSignlessIntOrFloat() ||
+          resultTy->getElementType().isa<ComplexType>())) {
+      return rewriter.notifyMatchFailure(
+          op, "mismatched operand/result types or iterator count");
+    }
+
+    // All-scalar pointwise ops inside of linalg ops are processes by
+    // ScalarHloToArithmeticPattern.
+    if (maxRank == 0 && isInBodyOfLinalgOps(op)) return failure();
+
+    // Find input/output values and types.
+    Value emptyTensor =
+        getEmptyTensorFor(rewriter, loc, *resultTy, op, adaptor.getOperands());
+
+    // Mapped inputs are cast to the same shape as the init tensor.
+    // Values from scalar inputs are extracted and used directly in the block.
+    SmallVector<Value> mappedInputs;
+    SmallVector<Value> scalarInputs;
+    for (Value input : adaptor.getOperands()) {
+      if (getRank(input) == maxRank) {
+        mappedInputs.push_back(coerceTensorShape(
+            rewriter, loc, cast<TypedValue<ShapedType>>(input),
+            emptyTensor.getType()));
+        scalarInputs.push_back(nullptr);
+      } else {
+        scalarInputs.push_back(rewriter.create<tensor::ExtractOp>(loc, input));
+      }
+    }
+
+    auto mapOp = rewriter.create<linalg::MapOp>(
+        loc, mappedInputs, emptyTensor,
+        [&](OpBuilder& b, Location loc, ValueRange args) {
+          Value innerResult = stablehlo::StableHloOpToStdScalarOp::mapOp(
+              op, getElementTypeOrSelf(emptyTensor),
+              interleaveScalarAndBlockArgs(scalarInputs, args), &b);
+
+          b.create<linalg::YieldOp>(loc, innerResult);
+        },
+        linalg::getPrunedAttributeList(op));
+
+    rewriter.replaceOp(op, mapOp->getResults());
+    return success();
+  }
+
+ protected:
+  int64_t getRank(Value v) const {
+    return v.getType().cast<ShapedType>().getRank();
+  }
+
+  int64_t getMaxRank(typename OpTy::Adaptor adaptor) const {
+    int64_t maxRank = 0;
+    for (auto operand : adaptor.getOperands()) {
+      maxRank = std::max(maxRank, getRank(operand));
+    }
+    return maxRank;
+  }
+
+  // Inserts block arguments in places where scalar inputs have a nullptr.
+  SmallVector<Value> interleaveScalarAndBlockArgs(ValueRange scalarInputs,
+                                                  ValueRange blockArgs) const {
+    SmallVector<Value> result;
+    auto argsIter = blockArgs.begin();
+    for (Value scalarInput : scalarInputs) {
+      if (scalarInput) {
+        result.push_back(scalarInput);
+      } else {
+        result.push_back(*argsIter);
+        ++argsIter;
+      }
+    }
+    return result;
+  }
+};
+
+class SetDimensionSizeConverter
+    : public OpConversionPattern<stablehlo::SetDimensionSizeOp> {
+ public:
+  using OpConversionPattern::OpConversionPattern;
+
+  LogicalResult matchAndRewrite(
+      stablehlo::SetDimensionSizeOp setDimensionSizeOp, OpAdaptor adaptor,
+      ConversionPatternRewriter& rewriter) const final {
+    // We can lower SetDimensionSize to tensor extract. This turns into a
+    // regular dynamic shape. Note that the bounds annotation is still around
+    // but may be no longer valid depending on choices made by bufferization.
+    Location loc = setDimensionSizeOp.getLoc();
+    auto resultType = setDimensionSizeOp.getType().cast<RankedTensorType>();
+
+    SmallVector<OpFoldResult> offsets(resultType.getRank(),
+                                      rewriter.getIndexAttr(0));
+    SmallVector<OpFoldResult> strides(resultType.getRank(),
+                                      rewriter.getIndexAttr(1));
+    SmallVector<OpFoldResult> sizes(llvm::map_range(
+        resultType.getShape(), [&](int64_t dim) -> OpFoldResult {
+          return rewriter.getIndexAttr(dim);
+        }));
+    Value dimensionSize =
+        rewriter.create<tensor::ExtractOp>(loc, setDimensionSizeOp.getSize());
+    sizes[setDimensionSizeOp.getDimension()] =
+        rewriter
+            .create<arith::IndexCastOp>(loc, rewriter.getIndexType(),
+                                        dimensionSize)
+            .getResult();
+
+    rewriter.replaceOpWithNewOp<tensor::ExtractSliceOp>(
+        setDimensionSizeOp, resultType, adaptor.getOperand(), offsets, sizes,
+        strides);
+    return success();
+  }
+};
+
+struct ConvertStableHloToLinalg
+    : impl::ConvertStableHloToLinalgBase<ConvertStableHloToLinalg> {
+  using ConvertStableHloToLinalgBase::ConvertStableHloToLinalgBase;
+
+  void getDependentDialects(DialectRegistry& registry) const override {
+    registry.insert<bufferization::BufferizationDialect, linalg::LinalgDialect,
+                    scf::SCFDialect, complex::ComplexDialect, math::MathDialect,
+                    memref::MemRefDialect, shape::ShapeDialect>();
+  }
+
+  void runOnOperation() override {
+    MLIRContext& ctx = getContext();
+    RewritePatternSet patterns(&ctx);
+    ConversionTarget target(ctx);
+    target.addLegalDialect<
+        bufferization::BufferizationDialect, arith::ArithDialect,
+        complex::ComplexDialect, linalg::LinalgDialect, math::MathDialect,
+        tensor::TensorDialect, sparse_tensor::SparseTensorDialect,
+        scf::SCFDialect, shape::ShapeDialect>();
+
+    target.addLegalOp<UnrealizedConversionCastOp>();
+
+    auto typeConverter = createStableHloToLinalgTypeConverter();
+    auto func = getOperation();
+    // TODO: Port `populateScalarHloToArithmeticConversionPatterns`.
+    populateStableHloToLinalgConversionPatterns(&ctx, *typeConverter, &patterns,
+                                                this->enablePrimitiveOps);
+    if (failed(applyPartialConversion(func, target, std::move(patterns)))) {
+      signalPassFailure();
+    }
+  }
+};
+
+}  // namespace
+
+void populateStableHloToLinalgConversionPatterns(MLIRContext* context,
+                                                 TypeConverter& typeConverter,
+                                                 RewritePatternSet* patterns,
+                                                 bool enablePrimitiveOps) {
+  // clang-format off
+  patterns->add<
+      BitcastConvertConverter,
+      ConcatenateConverter,
+      ConstConverterTensor,
+      EinsumToLinalgConverter,
+      RealDynamicSliceConverter,
+      ReshapeOpConverter,
+      ReverseConverter,
+      SetDimensionSizeConverter,
+      SliceConverter,
+      DynamicSliceConverter,
+      DynamicUpdateSliceConverter,
+      PadOpConversion,
+      PadOpNegativePaddingConversion,
+      TorchIndexSelectOpConversion,
+      SelectAndScatterNoOverlapConverter,
+      ReduceRegionReturnOpConversion>(typeConverter, context);
+
+  if (enablePrimitiveOps) {
+    patterns->add<
+      BroadcastInDimOpToBroadcastConverter,
+      BroadcastOpToBroadcastConverter,
+      DynamicBroadcastInDimOpToBroadcastConverter,
+      IotaToMapConverter<stablehlo::IotaOp>,
+      IotaToMapConverter<stablehlo::DynamicIotaOp>,
+      MapOpToMapConverter,
+      PointwiseToLinalgMapConverter<stablehlo::AbsOp>,
+      PointwiseToLinalgMapConverter<stablehlo::AddOp>,
+      PointwiseToLinalgMapConverter<stablehlo::AndOp>,
+      PointwiseToLinalgMapConverter<stablehlo::Atan2Op>,
+      PointwiseToLinalgMapConverter<stablehlo::BitcastConvertOp>,
+      PointwiseToLinalgMapConverter<stablehlo::CbrtOp>,
+      PointwiseToLinalgMapConverter<stablehlo::CeilOp>,
+      PointwiseToLinalgMapConverter<stablehlo::ClampOp>,
+      PointwiseToLinalgMapConverter<stablehlo::ClzOp>,
+      PointwiseToLinalgMapConverter<stablehlo::CompareOp>,
+      PointwiseToLinalgMapConverter<stablehlo::ComplexOp>,
+      PointwiseToLinalgMapConverter<stablehlo::ConvertOp>,
+      PointwiseToLinalgMapConverter<stablehlo::CosineOp>,
+      PointwiseToLinalgMapConverter<stablehlo::DivOp>,
+      PointwiseToLinalgMapConverter<stablehlo::ExpOp>,
+      PointwiseToLinalgMapConverter<stablehlo::Expm1Op>,
+      PointwiseToLinalgMapConverter<stablehlo::FloorOp>,
+      PointwiseToLinalgMapConverter<stablehlo::ImagOp>,
+      PointwiseToLinalgMapConverter<stablehlo::IsFiniteOp>,
+      PointwiseToLinalgMapConverter<stablehlo::Log1pOp>,
+      PointwiseToLinalgMapConverter<stablehlo::LogOp>,
+      PointwiseToLinalgMapConverter<stablehlo::LogisticOp>,
+      PointwiseToLinalgMapConverter<stablehlo::MaxOp>,
+      PointwiseToLinalgMapConverter<stablehlo::MinOp>,
+      PointwiseToLinalgMapConverter<stablehlo::MulOp>,
+      PointwiseToLinalgMapConverter<stablehlo::NegOp>,
+      PointwiseToLinalgMapConverter<stablehlo::NotOp>,
+      PointwiseToLinalgMapConverter<stablehlo::OrOp>,
+      PointwiseToLinalgMapConverter<stablehlo::PopulationCountOp>,
+      PointwiseToLinalgMapConverter<stablehlo::PowOp>,
+      PointwiseToLinalgMapConverter<stablehlo::RealOp>,
+      PointwiseToLinalgMapConverter<stablehlo::ReducePrecisionOp>,
+      PointwiseToLinalgMapConverter<stablehlo::RemOp>,
+      PointwiseToLinalgMapConverter<stablehlo::RoundNearestEvenOp>,
+      PointwiseToLinalgMapConverter<stablehlo::RoundOp>,
+      PointwiseToLinalgMapConverter<stablehlo::RsqrtOp>,
+      PointwiseToLinalgMapConverter<stablehlo::SelectOp>,
+      PointwiseToLinalgMapConverter<stablehlo::ShiftLeftOp>,
+      PointwiseToLinalgMapConverter<stablehlo::ShiftRightArithmeticOp>,
+      PointwiseToLinalgMapConverter<stablehlo::ShiftRightLogicalOp>,
+      PointwiseToLinalgMapConverter<stablehlo::SignOp>,
+      PointwiseToLinalgMapConverter<stablehlo::SineOp>,
+      PointwiseToLinalgMapConverter<stablehlo::SqrtOp>,
+      PointwiseToLinalgMapConverter<stablehlo::SubtractOp>,
+      PointwiseToLinalgMapConverter<stablehlo::TanhOp>,
+      PointwiseToLinalgMapConverter<stablehlo::XorOp>,
+      ReduceOpToReduceConverter,
+      TransposeOpToTransposeConverter
+    >(typeConverter, context);
+  } else {
+    patterns->add<
+      BroadcastConverter<stablehlo::BroadcastOp>,
+      IotaConverter<stablehlo::IotaOp>,
+      IotaConverter<stablehlo::DynamicIotaOp>,
+      HloBroadcastInDimConverter,
+      HloDynamicBroadcastInDimConverter,
+      MapOpToGenericConverter,
+      PointwiseToLinalgConverter<stablehlo::AbsOp>,
+      PointwiseToLinalgConverter<stablehlo::AddOp>,
+      PointwiseToLinalgConverter<stablehlo::AndOp>,
+      PointwiseToLinalgConverter<stablehlo::Atan2Op>,
+      PointwiseToLinalgConverter<stablehlo::BitcastConvertOp>,
+      PointwiseToLinalgConverter<stablehlo::CbrtOp>,
+      PointwiseToLinalgConverter<stablehlo::CeilOp>,
+      PointwiseToLinalgConverter<stablehlo::ClampOp>,
+      PointwiseToLinalgConverter<stablehlo::ClzOp>,
+      PointwiseToLinalgConverter<stablehlo::CompareOp>,
+      PointwiseToLinalgConverter<stablehlo::ComplexOp>,
+      PointwiseToLinalgConverter<stablehlo::ConvertOp>,
+      PointwiseToLinalgConverter<stablehlo::CosineOp>,
+      PointwiseToLinalgConverter<stablehlo::DivOp>,
+      PointwiseToLinalgConverter<stablehlo::ExpOp>,
+      PointwiseToLinalgConverter<stablehlo::Expm1Op>,
+      PointwiseToLinalgConverter<stablehlo::FloorOp>,
+      PointwiseToLinalgConverter<stablehlo::ImagOp>,
+      PointwiseToLinalgConverter<stablehlo::IsFiniteOp>,
+      PointwiseToLinalgConverter<stablehlo::Log1pOp>,
+      PointwiseToLinalgConverter<stablehlo::LogOp>,
+      PointwiseToLinalgConverter<stablehlo::LogisticOp>,
+      PointwiseToLinalgConverter<stablehlo::MaxOp>,
+      PointwiseToLinalgConverter<stablehlo::MinOp>,
+      PointwiseToLinalgConverter<stablehlo::MulOp>,
+      PointwiseToLinalgConverter<stablehlo::NegOp>,
+      PointwiseToLinalgConverter<stablehlo::NotOp>,
+      PointwiseToLinalgConverter<stablehlo::OrOp>,
+      PointwiseToLinalgConverter<stablehlo::PopulationCountOp>,
+      PointwiseToLinalgConverter<stablehlo::PowOp>,
+      PointwiseToLinalgConverter<stablehlo::RealOp>,
+      PointwiseToLinalgConverter<stablehlo::ReducePrecisionOp>,
+      PointwiseToLinalgConverter<stablehlo::RemOp>,
+      PointwiseToLinalgConverter<stablehlo::RoundNearestEvenOp>,
+      PointwiseToLinalgConverter<stablehlo::RoundOp>,
+      PointwiseToLinalgConverter<stablehlo::RsqrtOp>,
+      PointwiseToLinalgConverter<stablehlo::SelectOp>,
+      PointwiseToLinalgConverter<stablehlo::ShiftLeftOp>,
+      PointwiseToLinalgConverter<stablehlo::ShiftRightArithmeticOp>,
+      PointwiseToLinalgConverter<stablehlo::ShiftRightLogicalOp>,
+      PointwiseToLinalgConverter<stablehlo::SignOp>,
+      PointwiseToLinalgConverter<stablehlo::SineOp>,
+      PointwiseToLinalgConverter<stablehlo::SqrtOp>,
+      PointwiseToLinalgConverter<stablehlo::SubtractOp>,
+      PointwiseToLinalgConverter<stablehlo::TanhOp>,
+      PointwiseToLinalgConverter<stablehlo::XorOp>,
+      ReduceOpToGenericConverter,
+      TransposeConverter<stablehlo::TransposeOp>
+    >(typeConverter, context);
+  }
+
+  // Ensure specialized patterns are higher priority than their generic
+  // versions.
+  patterns->add<
+      DotOpConversion<DotOperationType::kMatrixMatrix, linalg::MatmulOp>,
+      DotOpConversion<DotOperationType::kMatrixVector, linalg::MatvecOp>,
+      DotOpConversion<DotOperationType::kVectorMatrix, linalg::VecmatOp>,
+      DotOpConversion<DotOperationType::kVectorDot, linalg::DotOp>,
+      DotGeneralBatchMatMulOpConversion>(typeConverter, context,
+                                         PatternBenefit(2));
+  patterns->add<
+      DotGeneralOpConversion>(typeConverter, context, PatternBenefit(1));
+  linalg::populateEraseUnusedOperandsAndResultsPatterns(*patterns);
+  // clang-format on
+}
+
+std::unique_ptr<TypeConverter> createStableHloToLinalgTypeConverter() {
+  return std::make_unique<LinalgTypeConverter>();
+}
+
+}  // namespace mlir::iree_compiler::stablehlo
diff --git a/compiler/src/iree/compiler/InputConversion/StableHLO/TypeConversion.cpp b/compiler/src/iree/compiler/InputConversion/StableHLO/TypeConversion.cpp
new file mode 100644
index 000000000000..832d25263f57
--- /dev/null
+++ b/compiler/src/iree/compiler/InputConversion/StableHLO/TypeConversion.cpp
@@ -0,0 +1,92 @@
+// Copyright 2021 The IREE Authors
+//
+// Licensed under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+#include "iree/compiler/InputConversion/StableHLO/TypeConversion.h"
+
+#include <optional>
+
+#include "mlir/Dialect/Func/IR/FuncOps.h"
+#include "mlir/Dialect/Func/Transforms/FuncConversions.h"
+#include "mlir/Dialect/Tensor/IR/Tensor.h"
+#include "mlir/IR/Builders.h"
+#include "mlir/IR/BuiltinOps.h"
+#include "mlir/IR/BuiltinTypes.h"
+#include "mlir/IR/TypeUtilities.h"
+#include "mlir/IR/Types.h"
+#include "mlir/IR/Value.h"
+#include "stablehlo/dialect/StablehloOps.h"
+
+namespace mlir::iree_compiler::stablehlo {
+
+namespace {
+
+Type convertInteger(IntegerType intType) {
+  return IntegerType::get(intType.getContext(),
+                          intType.getIntOrFloatBitWidth());
+}
+
+Type convertShapedType(ShapedType shapedType) {
+  if (auto intType = shapedType.getElementType().dyn_cast<IntegerType>())
+    return shapedType.clone(convertInteger(intType));
+  return shapedType;
+}
+
+std::optional<Value> materializeCastFromIllegal(OpBuilder& builder, Type type,
+                                                ValueRange inputs,
+                                                Location loc) {
+  Type fromType = getElementTypeOrSelf(inputs[0].getType());
+  Type toType = getElementTypeOrSelf(type);
+  if ((!fromType.isSignedInteger() && !fromType.isUnsignedInteger()) ||
+      !toType.isSignlessInteger())
+    return std::nullopt;
+  // Use unrealized conversion casts to do signful->signless conversions.
+  return builder.create<UnrealizedConversionCastOp>(loc, type, inputs[0])
+      ->getResult(0);
+}
+
+std::optional<Value> materializeCastToIllegal(OpBuilder& builder, Type type,
+                                              ValueRange inputs, Location loc) {
+  Type fromType = getElementTypeOrSelf(inputs[0].getType());
+  Type toType = getElementTypeOrSelf(type);
+  if (!fromType.isSignlessInteger() ||
+      (!toType.isSignedInteger() && !toType.isUnsignedInteger()))
+    return std::nullopt;
+  // Use unrealized conversion casts to do signless->signful conversions.
+  return builder.create<UnrealizedConversionCastOp>(loc, type, inputs[0])
+      ->getResult(0);
+}
+
+std::optional<Value> scalarToTensor(OpBuilder& builder, Type /*type*/,
+                                    ValueRange inputs, Location loc) {
+  assert(inputs.size() == 1);
+  if (inputs.front().getType().isa<ShapedType>()) {
+    return std::nullopt;
+  }
+  return builder
+      .create<tensor::FromElementsOp>(
+          loc, RankedTensorType::get({}, inputs.front().getType()),
+          inputs.front())
+      .getResult();
+}
+
+}  // namespace
+
+RemoveSignTypeConverter::RemoveSignTypeConverter() {
+  addConversion([](Type type) { return type; });
+
+  addConversion(convertInteger);
+  addConversion(convertShapedType);
+
+  addArgumentMaterialization(materializeCastFromIllegal);
+  addSourceMaterialization(materializeCastToIllegal);
+  addTargetMaterialization(materializeCastFromIllegal);
+}
+
+LinalgTypeConverter::LinalgTypeConverter() : RemoveSignTypeConverter() {
+  addArgumentMaterialization(scalarToTensor);
+}
+
+}  // namespace mlir::iree_compiler::stablehlo
diff --git a/compiler/src/iree/compiler/InputConversion/StableHLO/TypeConversion.h b/compiler/src/iree/compiler/InputConversion/StableHLO/TypeConversion.h
new file mode 100644
index 000000000000..47d249ba6a5c
--- /dev/null
+++ b/compiler/src/iree/compiler/InputConversion/StableHLO/TypeConversion.h
@@ -0,0 +1,34 @@
+// Copyright 2021 The IREE Authors
+//
+// Licensed under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+#ifndef IREE_COMPILER_INPUTCONVERSION_STABLEHLO_TYPE_CONVERSION_H
+#define IREE_COMPILER_INPUTCONVERSION_STABLEHLO_TYPE_CONVERSION_H
+
+#include "mlir/IR/Attributes.h"
+#include "mlir/IR/Dialect.h"
+#include "mlir/Transforms/DialectConversion.h"
+
+namespace mlir::iree_compiler::stablehlo {
+
+// Type converter to use as part of lowerings from dialects that carry signs
+// in their types to those that are signless.
+class RemoveSignTypeConverter : public TypeConverter {
+ public:
+  RemoveSignTypeConverter();
+};
+
+// Type converter which adds additional materializations (beyond signless)
+// that are needed as part of the HloToLinalg conversion patterns.
+// This is the type converter used by the test pass and is the sanctioned
+// way to use the underlying patterns.
+class LinalgTypeConverter : public RemoveSignTypeConverter {
+ public:
+  LinalgTypeConverter();
+};
+
+}  // namespace mlir::iree_compiler::stablehlo
+
+#endif  // IREE_COMPILER_INPUTCONVERSION_STABLEHLO_TYPE_CONVERSION_H
diff --git a/compiler/src/iree/compiler/InputConversion/StableHLO/test/BUILD.bazel b/compiler/src/iree/compiler/InputConversion/StableHLO/test/BUILD.bazel
new file mode 100644
index 000000000000..68394d284d90
--- /dev/null
+++ b/compiler/src/iree/compiler/InputConversion/StableHLO/test/BUILD.bazel
@@ -0,0 +1,30 @@
+# Copyright 2023 The IREE Authors
+#
+# Licensed under the Apache License v2.0 with LLVM Exceptions.
+# See https://llvm.org/LICENSE.txt for license information.
+# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+# Tests for common transforms.
+
+load("//build_tools/bazel:iree_lit_test.bzl", "iree_lit_test_suite")
+load("//build_tools/bazel:enforce_glob.bzl", "enforce_glob")
+
+package(
+    features = ["layering_check"],
+    licenses = ["notice"],  # Apache 2.0
+)
+
+iree_lit_test_suite(
+    name = "lit",
+    srcs = enforce_glob(
+        [
+            "stablehlo_to_linalg.mlir",
+        ],
+        include = ["*.mlir"],
+    ),
+    cfg = "//compiler:lit.cfg.py",
+    tools = [
+        "//tools:iree-opt",
+        "@llvm-project//llvm:FileCheck",
+    ],
+)
diff --git a/compiler/src/iree/compiler/InputConversion/StableHLO/test/CMakeLists.txt b/compiler/src/iree/compiler/InputConversion/StableHLO/test/CMakeLists.txt
new file mode 100644
index 000000000000..2e0bead2a227
--- /dev/null
+++ b/compiler/src/iree/compiler/InputConversion/StableHLO/test/CMakeLists.txt
@@ -0,0 +1,23 @@
+################################################################################
+# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from           #
+# compiler/src/iree/compiler/InputConversion/StableHLO/test/BUILD.bazel        #
+#                                                                              #
+# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary   #
+# CMake-only content.                                                          #
+#                                                                              #
+# To disable autogeneration for this file entirely, delete this header.        #
+################################################################################
+
+iree_add_all_subdirs()
+
+iree_lit_test_suite(
+  NAME
+    lit
+  SRCS
+    "stablehlo_to_linalg.mlir"
+  TOOLS
+    FileCheck
+    iree-opt
+)
+
+### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ###
diff --git a/compiler/src/iree/compiler/InputConversion/StableHLO/test/stablehlo_to_linalg.mlir b/compiler/src/iree/compiler/InputConversion/StableHLO/test/stablehlo_to_linalg.mlir
new file mode 100644
index 000000000000..1da754c79267
--- /dev/null
+++ b/compiler/src/iree/compiler/InputConversion/StableHLO/test/stablehlo_to_linalg.mlir
@@ -0,0 +1,2054 @@
+// RUN: iree-opt %s --iree-stablehlo-to-linalg --split-input-file \
+// RUN:   --canonicalize | FileCheck --enable-var-scope=false %s
+
+// RUN: iree-opt %s --iree-stablehlo-to-linalg="enable-primitive-ops=true" \
+// RUN:   --split-input-file --canonicalize | \
+// RUN:   FileCheck %s --enable-var-scope=false --check-prefix=CHECK-PRIMITIVE
+
+// CHECK: #map = affine_map<(d0, d1) -> (d0, d1)>
+// CHECK-LABEL: func @float_add
+// CHECK-PRIMITIVE-LABEL: func @float_add
+func.func @float_add(%lhs: tensor<2x2xf32>,
+                %rhs: tensor<2x2xf32>) -> tensor<2x2xf32> {
+  // CHECK: linalg.generic
+  // CHECK-SAME: {someattr}
+  // CHECK: ^{{[a-z0-9_]*}}
+  // CHECK-SAME: %[[ARG0:[a-zA-Z0-9_]*]]: f32
+  // CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]*]]: f32
+  // CHECK: %[[RESULT:[a-zA-Z0-9_]*]] = arith.addf %[[ARG0]], %[[ARG1]]
+  // CHECK: linalg.yield %[[RESULT]]
+
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: arith.addf
+  %0 = "stablehlo.add"(%lhs, %rhs) {someattr}
+      : (tensor<2x2xf32>, tensor<2x2xf32>) -> tensor<2x2xf32>
+  func.return %0 : tensor<2x2xf32>
+}
+
+// -----
+
+// CHECK-LABEL: func @float_add_dynamic_encoding
+// CHECK-PRIMITIVE-LABEL: func @float_add_dynamic_encoding
+func.func @float_add_dynamic_encoding(
+  %lhs: tensor<2x?xf32, #stablehlo.type_extensions<bounds = [?, 2]>>,
+  %rhs: tensor<2x?xf32, #stablehlo.type_extensions<bounds = [?, 2]>>)
+    -> tensor<2x?xf32, #stablehlo.type_extensions<bounds = [?, 2]>> {
+  // CHECK: linalg.generic
+  // CHECK: arith.addf
+  // CHECK: linalg.yield
+
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: arith.addf
+  %0 = "stablehlo.add"(%lhs, %rhs)
+      : (tensor<2x?xf32, #stablehlo.type_extensions<bounds = [?, 2]>>,
+         tensor<2x?xf32, #stablehlo.type_extensions<bounds = [?, 2]>>)
+      -> tensor<2x?xf32, #stablehlo.type_extensions<bounds = [?, 2]>>
+  func.return %0 : tensor<2x?xf32, #stablehlo.type_extensions<bounds = [?, 2]>>
+}
+
+// -----
+
+// CHECK-LABEL: integer_add
+// CHECK-PRIMITIVE-LABEL: integer_add
+func.func @integer_add(%lhs: tensor<2x2xi32>,
+                  %rhs: tensor<2x2xi32>) -> tensor<2x2xi32> {
+  // CHECK: linalg.generic
+  // CHECK: addi
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: addi
+  %0 = "stablehlo.add"(%lhs, %rhs) : (tensor<2x2xi32>,
+                                    tensor<2x2xi32>) -> tensor<2x2xi32>
+  func.return %0 : tensor<2x2xi32>
+}
+
+// -----
+
+// CHECK-LABEL: complex_add
+// CHECK-PRIMITIVE-LABEL: complex_add
+func.func @complex_add(%lhs: tensor<2x2xcomplex<f32>>,
+                  %rhs: tensor<2x2xcomplex<f32>>) -> tensor<2x2xcomplex<f32>> {
+  // CHECK: linalg.generic
+  // CHECK: complex.add
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: complex.add
+  %0 = "stablehlo.add"(%lhs, %rhs) : (tensor<2x2xcomplex<f32>>,
+      tensor<2x2xcomplex<f32>>) -> tensor<2x2xcomplex<f32>>
+  func.return %0 : tensor<2x2xcomplex<f32>>
+}
+
+// -----
+
+// CHECK-LABEL: func @complex_atan2
+// CHECK-PRIMITIVE-LABEL: func @complex_atan2
+func.func @complex_atan2(%lhs: tensor<2x2xcomplex<f32>>,
+    %rhs: tensor<2x2xcomplex<f32>>) -> tensor<2x2xcomplex<f32>> {
+  %tensor_result = "stablehlo.atan2"(%lhs, %rhs)
+      : (tensor<2x2xcomplex<f32>>, tensor<2x2xcomplex<f32>>)
+      -> tensor<2x2xcomplex<f32>>
+  // CHECK: linalg.generic
+  // CHECK: complex.atan2
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: complex.atan2
+  func.return %tensor_result : tensor<2x2xcomplex<f32>>
+}
+
+
+// -----
+
+// CHECK-LABEL: func @float_mul
+// CHECK-PRIMITIVE-LABEL: func @float_mul
+func.func @float_mul(%lhs: tensor<2x2xf32>,
+                %rhs: tensor<2x2xf32>) -> tensor<2x2xf32> {
+  // CHECK: linalg.generic
+  // CHECK: mulf
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: mulf
+  %0 = "stablehlo.multiply"(%lhs, %rhs) : (tensor<2x2xf32>,
+                                    tensor<2x2xf32>) -> tensor<2x2xf32>
+  func.return %0 : tensor<2x2xf32>
+}
+
+// -----
+
+// CHECK-LABEL: func @integer_mul
+// CHECK-PRIMITIVE-LABEL: func @integer_mul
+func.func @integer_mul(%lhs: tensor<2x2xi32>,
+                  %rhs: tensor<2x2xi32>) -> tensor<2x2xi32> {
+  // CHECK: linalg.generic
+  // CHECK: muli
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: muli
+  %0 = "stablehlo.multiply"(%lhs, %rhs) : (tensor<2x2xi32>,
+                                    tensor<2x2xi32>) -> tensor<2x2xi32>
+  func.return %0 : tensor<2x2xi32>
+}
+
+// -----
+
+// CHECK-LABEL: func @complex_mul
+// CHECK-PRIMITIVE-LABEL: func @complex_mul
+func.func @complex_mul(%lhs: tensor<2x2xcomplex<f32>>,
+                  %rhs: tensor<2x2xcomplex<f32>>) -> tensor<2x2xcomplex<f32>> {
+  // CHECK: linalg.generic
+  // CHECK: complex.mul
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: complex.mul
+  %0 = "stablehlo.multiply"(%lhs, %rhs)
+          : (tensor<2x2xcomplex<f32>>, tensor<2x2xcomplex<f32>>)
+          -> tensor<2x2xcomplex<f32>>
+  func.return %0 : tensor<2x2xcomplex<f32>>
+}
+
+// -----
+
+// CHECK-LABEL: func @float_remainder
+// CHECK-PRIMITIVE-LABEL: func @float_remainder
+func.func @float_remainder(%lhs: tensor<2x2xf32>,
+                      %rhs: tensor<2x2xf32>) -> tensor<2x2xf32> {
+  // CHECK: linalg.generic
+  // CHECK: remf
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: remf
+  %0 = "stablehlo.remainder"(%lhs, %rhs) : (tensor<2x2xf32>,
+                                    tensor<2x2xf32>) -> tensor<2x2xf32>
+  func.return %0 : tensor<2x2xf32>
+}
+
+// -----
+
+// CHECK-LABEL: func @integer_remainder
+// CHECK-PRIMITIVE-LABEL: func @integer_remainder
+func.func @integer_remainder(%lhs: tensor<2x2xi32>,
+                        %rhs: tensor<2x2xi32>) -> tensor<2x2xi32> {
+  // CHECK: linalg.generic
+  // CHECK: arith.remsi
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: arith.remsi
+  %0 = "stablehlo.remainder"(%lhs, %rhs) : (tensor<2x2xi32>,
+                                          tensor<2x2xi32>) -> tensor<2x2xi32>
+  func.return %0 : tensor<2x2xi32>
+}
+
+// -----
+
+// CHECK-LABEL: func @population_count_integer
+// CHECK-PRIMITIVE-LABEL: func @population_count_integer
+func.func @population_count_integer(%lhs: tensor<2x2xi32>) -> tensor<2x2xi32> {
+  // CHECK: linalg.generic
+  // CHECK: math.ctpop
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: math.ctpop
+  %0 = "stablehlo.popcnt"(%lhs) : (tensor<2x2xi32>) -> tensor<2x2xi32>
+  func.return %0 : tensor<2x2xi32>
+}
+
+// -----
+
+// CHECK-LABEL: func @complex_sqrt
+// CHECK-PRIMITIVE-LABEL: func @complex_sqrt
+func.func @complex_sqrt(%operand: tensor<2x2xcomplex<f32>>) -> tensor<2x2xcomplex<f32>> {
+  %tensor_result = "stablehlo.sqrt"(%operand)
+      : (tensor<2x2xcomplex<f32>>) -> tensor<2x2xcomplex<f32>>
+  // CHECK: linalg.generic
+  // CHECK: complex.sqrt
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: complex.sqrt
+  func.return %tensor_result : tensor<2x2xcomplex<f32>>
+}
+
+// -----
+
+// CHECK-LABEL: func @float_rsqrt
+// CHECK-PRIMITIVE-LABEL: func @float_rsqrt
+func.func @float_rsqrt(%operand: tensor<2x2xf32>) -> tensor<2x2xf32> {
+  %tensor_result = "stablehlo.rsqrt"(%operand)
+      : (tensor<2x2xf32>) -> tensor<2x2xf32>
+  // CHECK: linalg.generic
+  // CHECK: rsqrt
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: rsqrt
+  func.return %tensor_result : tensor<2x2xf32>
+}
+
+// -----
+
+// CHECK-LABEL: func @complex_rsqrt
+// CHECK-PRIMITIVE-LABEL: func @complex_rsqrt
+func.func @complex_rsqrt(%operand: tensor<2x2xcomplex<f32>>)
+    -> tensor<2x2xcomplex<f32>> {
+  %tensor_result = "stablehlo.rsqrt"(%operand)
+      : (tensor<2x2xcomplex<f32>>) -> tensor<2x2xcomplex<f32>>
+  // CHECK: linalg.generic
+  // CHECK: complex.rsqrt
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: complex.rsqrt
+  func.return %tensor_result : tensor<2x2xcomplex<f32>>
+}
+
+// -----
+
+// CHECK-LABEL: func @float_cbrt
+// CHECK-PRIMITIVE-LABEL: func @float_cbrt
+func.func @float_cbrt(%operand: tensor<2x2xf32>) -> tensor<2x2xf32> {
+  %tensor_result = "stablehlo.cbrt"(%operand)
+      : (tensor<2x2xf32>) -> tensor<2x2xf32>
+  // CHECK: ^{{[a-z0-9_]*}}
+  // CHECK-SAME: %[[IN:[a-zA-Z0-9_]*]]: f32
+  // CHECK: %[[RESULT:.+]] = math.cbrt %[[IN]]
+  // CHECK: linalg.yield %[[RESULT]]
+
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: math.cbrt
+  func.return %tensor_result : tensor<2x2xf32>
+}
+
+// -----
+
+
+// CHECK-LABEL: func @float_sub
+// CHECK-PRIMITIVE-LABEL: func @float_sub
+func.func @float_sub(%lhs: tensor<2x2xf32>,
+                %rhs: tensor<2x2xf32>) -> tensor<2x2xf32> {
+  // CHECK: linalg.generic
+  // CHECK: subf
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: subf
+  %0 = "stablehlo.subtract"(%lhs, %rhs) : (tensor<2x2xf32>,
+                                    tensor<2x2xf32>) -> tensor<2x2xf32>
+  func.return %0 : tensor<2x2xf32>
+}
+
+// -----
+
+// CHECK-LABEL: func @integer_sub
+// CHECK-PRIMITIVE-LABEL: func @integer_sub
+func.func @integer_sub(%lhs: tensor<2x2xi32>,
+                  %rhs: tensor<2x2xi32>) -> tensor<2x2xi32> {
+  // CHECK: linalg.generic
+  // CHECK: subi
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: subi
+  %0 = "stablehlo.subtract"(%lhs, %rhs) : (tensor<2x2xi32>,
+                                    tensor<2x2xi32>) -> tensor<2x2xi32>
+  func.return %0 : tensor<2x2xi32>
+}
+
+// -----
+
+// CHECK-LABEL: complex_sub
+// CHECK-PRIMITIVE-LABEL: complex_sub
+func.func @complex_sub(%lhs: tensor<2x2xcomplex<f32>>,
+                  %rhs: tensor<2x2xcomplex<f32>>) -> tensor<2x2xcomplex<f32>> {
+  // CHECK: linalg.generic
+  // CHECK: complex.sub
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: complex.sub
+  %0 = "stablehlo.subtract"(%lhs, %rhs) : (tensor<2x2xcomplex<f32>>,
+      tensor<2x2xcomplex<f32>>) -> tensor<2x2xcomplex<f32>>
+  func.return %0 : tensor<2x2xcomplex<f32>>
+}
+
+// -----
+
+// CHECK-LABEL: func @float_abs
+// CHECK-PRIMITIVE-LABEL: func @float_abs
+func.func @float_abs(%arg0: tensor<2x2xf32>) -> tensor<2x2xf32> {
+  // CHECK: linalg.generic
+  // CHECK-SAME: {someattr}
+  // CHECK: math.absf
+  // CHECK-PRIMITIVE: linalg.map { math.absf }
+  // CHECK-PRIMITIVE-SAME: ins(
+  // CHECK-PRIMITIVE-SAME: outs(
+  // CHECK-PRIMITIVE-SAME: {someattr}
+  %0 = "stablehlo.abs"(%arg0) {someattr} : (tensor<2x2xf32>) -> tensor<2x2xf32>
+  func.return %0 : tensor<2x2xf32>
+}
+
+// -----
+
+// CHECK-LABEL: func @float_exp
+// CHECK-PRIMITIVE-LABEL: func @float_exp
+func.func @float_exp(%arg0: tensor<2x2xf32>) -> tensor<2x2xf32> {
+  // CHECK: linalg.generic
+  // CHECK: exp
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: exp
+  %0 = "stablehlo.exponential"(%arg0) : (tensor<2x2xf32>) -> tensor<2x2xf32>
+  func.return %0 : tensor<2x2xf32>
+}
+
+// -----
+
+// CHECK-LABEL: func @complex_exp
+func.func @complex_exp(%arg0: tensor<2x2xcomplex<f32>>) -> tensor<2x2xcomplex<f32>> {
+  // CHECK: linalg.generic
+  // CHECK: complex.exp
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: complex.exp
+  %0 = "stablehlo.exponential"(%arg0) : (tensor<2x2xcomplex<f32>>)
+                                 -> tensor<2x2xcomplex<f32>>
+  func.return %0 : tensor<2x2xcomplex<f32>>
+}
+
+// -----
+
+// CHECK-LABEL: func @float_expm1
+func.func @float_expm1(%arg0: tensor<2x2xf32>) -> tensor<2x2xf32> {
+  // CHECK: linalg.generic
+  // CHECK: expm1
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: expm1
+  %0 = "stablehlo.exponential_minus_one"(%arg0) : (tensor<2x2xf32>) -> tensor<2x2xf32>
+  func.return %0 : tensor<2x2xf32>
+}
+
+// -----
+
+// CHECK-LABEL: func @complex_expm1
+func.func @complex_expm1(%arg0: tensor<2x2xcomplex<f32>>)
+    -> tensor<2x2xcomplex<f32>> {
+  // CHECK: linalg.generic
+  // CHECK: complex.expm1
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: complex.expm1
+  %0 = "stablehlo.exponential_minus_one"(%arg0)
+    : (tensor<2x2xcomplex<f32>>) -> tensor<2x2xcomplex<f32>>
+  func.return %0 : tensor<2x2xcomplex<f32>>
+}
+
+// -----
+
+// CHECK-LABEL: func @float_log
+func.func @float_log(%arg0: tensor<2x2xf32>) -> tensor<2x2xf32> {
+  // CHECK: linalg.generic
+  // CHECK: math.log
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: math.log
+  %0 = "stablehlo.log"(%arg0) : (tensor<2x2xf32>) -> tensor<2x2xf32>
+  func.return %0 : tensor<2x2xf32>
+}
+
+// -----
+
+// CHECK-LABEL: func @complex_log
+func.func @complex_log(%arg0: tensor<2x2xcomplex<f32>>) -> tensor<2x2xcomplex<f32>> {
+  // CHECK: linalg.generic
+  // CHECK: complex.log
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: complex.log
+  %0 = "stablehlo.log"(%arg0) : (tensor<2x2xcomplex<f32>>)
+                         -> tensor<2x2xcomplex<f32>>
+  func.return %0 : tensor<2x2xcomplex<f32>>
+}
+
+// -----
+
+// CHECK-LABEL: func @float_log1p
+// CHECK-PRIMITIVE-LABEL: func @float_log1p
+func.func @float_log1p(%arg0: tensor<2x2xf32>) -> tensor<2x2xf32> {
+  // CHECK: linalg.generic
+  // CHECK: math.log1p
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: math.log1p
+  %0 = "stablehlo.log_plus_one"(%arg0) : (tensor<2x2xf32>) -> tensor<2x2xf32>
+  func.return %0 : tensor<2x2xf32>
+}
+
+// -----
+
+// CHECK-LABEL: func @complex_log1p
+// CHECK-PRIMITIVE-LABEL: func @complex_log1p
+func.func @complex_log1p(%arg0: tensor<2x2xcomplex<f32>>)
+    -> tensor<2x2xcomplex<f32>> {
+  // CHECK: linalg.generic
+  // CHECK: complex.log1p
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: complex.log1p
+  %0 = "stablehlo.log_plus_one"(%arg0) : (tensor<2x2xcomplex<f32>>)
+                                  -> tensor<2x2xcomplex<f32>>
+  func.return %0 : tensor<2x2xcomplex<f32>>
+}
+
+// -----
+
+// CHECK-LABEL: func @float_logistic
+// CHECK-PRIMITIVE-LABEL: func @float_logistic
+func.func @float_logistic(%arg0: tensor<2x2xf32>) -> tensor<2x2xf32> {
+  // CHECK: %[[C1:.*]] = arith.constant 1.{{.*}}e+00
+  // CHECK: linalg.generic
+  // CHECK: ^bb0(%[[ARG:.*]]: f32, %{{.*}}: f32):
+  // CHECK: %[[NEG_ARG:.*]] = arith.negf %[[ARG]]
+  // CHECK: %[[EXP_NEG_ARG:.*]] = math.exp %[[NEG_ARG]]
+  // CHECK: %[[ONE_ADD_EXP_NEG_ARG:.*]] = arith.addf %[[EXP_NEG_ARG]], %[[C1]]
+  // CHECK: %[[RESULT:.*]] = arith.divf %[[C1]], %[[ONE_ADD_EXP_NEG_ARG]]
+  // CHECK: linalg.yield %[[RESULT]]
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: arith.negf
+  // CHECK-PRIMITIVE: math.exp
+  // CHECK-PRIMITIVE: arith.addf
+  // CHECK-PRIMITIVE: arith.divf
+  %0 = "stablehlo.logistic"(%arg0) : (tensor<2x2xf32>) -> tensor<2x2xf32>
+  func.return %0 : tensor<2x2xf32>
+}
+
+// -----
+
+// CHECK-LABEL: func @complex_logistic
+func.func @complex_logistic(%arg0: tensor<2x2xcomplex<f32>>) -> tensor<2x2xcomplex<f32>> {
+  // CHECK-DAG: %[[C0:.*]] = arith.constant 0
+  // CHECK-DAG: %[[C1:.*]] = arith.constant 1
+  // CHECK: linalg.generic
+  // CHECK: ^bb0(%[[ARG:.*]]: complex<f32>, %{{.*}}: complex<f32>):
+  // CHECK: %[[NEG_ARG:.*]] = complex.neg %[[ARG]]
+  // CHECK: %[[EXP_NEG_ARG:.*]] = complex.exp %[[NEG_ARG]]
+  // CHECK: %[[CC1:.*]] = complex.create %[[C1]], %[[C0]] : complex<f32>
+  // CHECK: %[[ONE_ADD_EXP_NEG_ARG:.*]] = complex.add %[[EXP_NEG_ARG]], %[[CC1]]
+  // CHECK: %[[RESULT:.*]] = complex.div %[[CC1]], %[[ONE_ADD_EXP_NEG_ARG]]
+  // CHECK: linalg.yield %[[RESULT]]
+  %0 = "stablehlo.logistic"(%arg0) : (tensor<2x2xcomplex<f32>>) -> tensor<2x2xcomplex<f32>>
+  func.return %0 : tensor<2x2xcomplex<f32>>
+}
+
+// -----
+
+// CHECK-LABEL: func @float_ceil
+// CHECK-PRIMITIVE-LABEL: func @float_ceil
+func.func @float_ceil(%arg0: tensor<2x2xf32>) -> tensor<2x2xf32> {
+  // CHECK: linalg.generic
+  // CHECK: math.ceil
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: math.ceil
+  %0 = "stablehlo.ceil"(%arg0) : (tensor<2x2xf32>) -> tensor<2x2xf32>
+  func.return %0 : tensor<2x2xf32>
+}
+
+// -----
+
+// CHECK-LABEL: func @floor
+// CHECK-PRIMITIVE-LABEL: func @floor
+func.func @floor(%input: tensor<2x2xf32>) -> tensor<2x2xf32> {
+  // CHECK: linalg.generic
+  // CHECK: math.floor
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: math.floor
+  %0 = "stablehlo.floor"(%input) : (tensor<2x2xf32>) -> tensor<2x2xf32>
+  func.return %0 : tensor<2x2xf32>
+}
+
+// -----
+
+// CHECK-LABEL: func @float_neg
+// CHECK-PRIMITIVE-LABEL: func @float_neg
+func.func @float_neg(%arg0: tensor<2x2xf32>) -> tensor<2x2xf32> {
+  // CHECK: linalg.generic
+  // CHECK: negf
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: negf
+  %0 = "stablehlo.negate"(%arg0) : (tensor<2x2xf32>) -> tensor<2x2xf32>
+  func.return %0 : tensor<2x2xf32>
+}
+
+// -----
+
+// CHECK-LABEL: func @complex_neg
+// CHECK-PRIMITIVE-LABEL: func @complex_neg
+func.func @complex_neg(%arg0: tensor<2x2xcomplex<f32>>) -> tensor<2x2xcomplex<f32>> {
+  // CHECK: linalg.generic
+  // CHECK: complex.neg
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: complex.neg
+  %0 = "stablehlo.negate"(%arg0) : (tensor<2x2xcomplex<f32>>)
+                            -> tensor<2x2xcomplex<f32>>
+  func.return %0 : tensor<2x2xcomplex<f32>>
+}
+
+// -----
+
+// CHECK-LABEL: func @complex_sign
+// CHECK-PRIMITIVE-LABEL: func @complex_sign
+func.func @complex_sign(
+    %arg0: tensor<2x2xcomplex<f32>>) -> tensor<2x2xcomplex<f32>> {
+  // CHECK: linalg.generic
+  // CHECK: complex.sign
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: complex.sign
+  %0 = "stablehlo.sign"(%arg0) : (tensor<2x2xcomplex<f32>>)
+                          -> tensor<2x2xcomplex<f32>>
+  func.return %0 : tensor<2x2xcomplex<f32>>
+}
+
+// -----
+
+// CHECK-LABEL: func @float_tanh
+// CHECK-PRIMITIVE-LABEL: func @float_tanh
+func.func @float_tanh(%arg0: tensor<2x2xf32>) -> tensor<2x2xf32> {
+  // CHECK: linalg.generic
+  // CHECK: tanh
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: tanh
+  %0 = "stablehlo.tanh"(%arg0) : (tensor<2x2xf32>) -> tensor<2x2xf32>
+  func.return %0 : tensor<2x2xf32>
+}
+
+// -----
+
+// CHECK-LABEL: func @complex_tanh
+// CHECK-PRIMITIVE-LABEL: func @complex_tanh
+func.func @complex_tanh(%operand: tensor<2x2xcomplex<f32>>) -> tensor<2x2xcomplex<f32>> {
+  %tensor_result = "stablehlo.tanh"(%operand)
+      : (tensor<2x2xcomplex<f32>>) -> tensor<2x2xcomplex<f32>>
+  // CHECK: linalg.generic
+  // CHECK: complex.tanh
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: complex.tanh
+  func.return %tensor_result : tensor<2x2xcomplex<f32>>
+}
+
+// -----
+
+// CHECK-LABEL: func @integer_and
+// CHECK-PRIMITIVE-LABEL: func @integer_and
+func.func @integer_and(%lhs: tensor<2x2xi32>,
+                  %rhs: tensor<2x2xi32>) -> tensor<2x2xi32> {
+  // CHECK: linalg.generic
+  // CHECK: and
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: and
+  %0 = "stablehlo.and"(%lhs, %rhs) : (tensor<2x2xi32>,
+                                    tensor<2x2xi32>) -> tensor<2x2xi32>
+  func.return %0 : tensor<2x2xi32>
+}
+
+// -----
+
+// CHECK-LABEL: func @integer_or
+// CHECK-PRIMITIVE-LABEL: func @integer_or
+func.func @integer_or(%lhs: tensor<2x2xi32>,
+                  %rhs: tensor<2x2xi32>) -> tensor<2x2xi32> {
+  // CHECK: linalg.generic
+  // CHECK: or
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: or
+  %0 = "stablehlo.or"(%lhs, %rhs) : (tensor<2x2xi32>,
+                                    tensor<2x2xi32>) -> tensor<2x2xi32>
+  func.return %0 : tensor<2x2xi32>
+}
+
+// -----
+
+// CHECK-LABEL: func @integer_xor
+// CHECK-PRIMITIVE-LABEL: func @integer_xor
+func.func @integer_xor(%lhs: tensor<2x2xi32>,
+                  %rhs: tensor<2x2xi32>) -> tensor<2x2xi32> {
+  // CHECK: linalg.generic
+  // CHECK: xor
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: xor
+  %0 = "stablehlo.xor"(%lhs, %rhs) : (tensor<2x2xi32>,
+                                    tensor<2x2xi32>) -> tensor<2x2xi32>
+  func.return %0 : tensor<2x2xi32>
+}
+
+// -----
+
+// CHECK-LABEL: func @count_leading_zeros
+// CHECK-PRIMITIVE-LABEL: func @count_leading_zeros
+func.func @count_leading_zeros(%lhs: tensor<2x2xi32>) -> tensor<2x2xi32> {
+  // CHECK: linalg.generic
+  // CHECK: math.ctlz
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: math.ctlz
+  %0 = "stablehlo.count_leading_zeros"(%lhs) : (tensor<2x2xi32>) -> tensor<2x2xi32>
+  func.return %0 : tensor<2x2xi32>
+}
+
+// -----
+
+// CHECK-LABEL: func @float_cmp
+// CHECK-PRIMITIVE-LABEL: func @float_cmp
+func.func @float_cmp(%lhs: tensor<2x2xf32>,
+                %rhs: tensor<2x2xf32>) -> (tensor<2x2xi1>) {
+  %0 = "stablehlo.compare"(%lhs, %rhs) {comparison_direction = #stablehlo<comparison_direction EQ>}
+          : (tensor<2x2xf32>, tensor<2x2xf32>) -> tensor<2x2xi1>
+  func.return %0 : tensor<2x2xi1>
+}
+// CHECK: tensor.empty() : tensor<2x2xi1>
+// CHECK: linalg.generic
+// CHECK-NEXT: ^bb0(%[[LHS_IN:.*]]: f32, %[[RHS_IN:.*]]: f32, %{{.*}}: i1):
+// CHECK-NEXT:   %[[RESULT:.*]] = arith.cmpf oeq, %[[LHS_IN]], %[[RHS_IN]] : f32
+// CHECK-NEXT:   linalg.yield %[[RESULT]] : i1
+// CHECK-PRIMITIVE: linalg.map
+// CHECK-PRIMITIVE: arith.cmpf
+
+// -----
+
+// CHECK-LABEL: func @float_cmp_ne
+// CHECK-PRIMITIVE-LABEL: func @float_cmp_ne
+func.func @float_cmp_ne(%lhs: tensor<2x2xf32>,
+                %rhs: tensor<2x2xf32>) -> (tensor<2x2xi1>) {
+  %0 = "stablehlo.compare"(%lhs, %rhs) {comparison_direction = #stablehlo<comparison_direction NE>}
+          : (tensor<2x2xf32>, tensor<2x2xf32>) -> tensor<2x2xi1>
+  func.return %0 : tensor<2x2xi1>
+}
+// CHECK: tensor.empty() : tensor<2x2xi1>
+// CHECK: linalg.generic
+// CHECK-NEXT: ^bb0(%[[LHS_IN:.*]]: f32, %[[RHS_IN:.*]]: f32, %{{.*}}: i1):
+// CHECK-NEXT:   %[[RESULT:.*]] = arith.cmpf une, %[[LHS_IN]], %[[RHS_IN]] : f32
+// CHECK-NEXT:   linalg.yield %[[RESULT]] : i1
+// CHECK-PRIMITIVE: linalg.map
+// CHECK-PRIMITIVE: arith.cmpf
+
+// -----
+
+// CHECK-LABEL: func @float_cmp_totalorder
+// CHECK-PRIMITIVE-LABEL: func @float_cmp_totalorder
+func.func @float_cmp_totalorder(%lhs: tensor<2x2xbf16>,
+                %rhs: tensor<2x2xbf16>) -> (tensor<2x2xi1>) {
+  %0 = "stablehlo.compare"(%lhs, %rhs) {
+    comparison_direction = #stablehlo<comparison_direction LT>,
+    compare_type = #stablehlo<comparison_type TOTALORDER>
+  } : (tensor<2x2xbf16>, tensor<2x2xbf16>) -> tensor<2x2xi1>
+  func.return %0 : tensor<2x2xi1>
+}
+// CHECK-DAG: %[[C0:.*]] = arith.constant 0 : i16
+// CHECK-DAG: %[[C32767:.*]] = arith.constant 32767 : i16
+// CHECK: tensor.empty() : tensor<2x2xi1>
+// CHECK: linalg.generic
+// CHECK-NEXT: ^bb0(%[[LHS_IN:.*]]: bf16, %[[RHS_IN:.*]]: bf16, %{{.*}}: i1):
+// CHECK-NEXT:   %[[LHS_INT:.*]] = arith.bitcast %[[LHS_IN]] : bf16 to i16
+// CHECK-NEXT:   %[[LHS_CMP:.*]] = arith.cmpi slt, %[[LHS_INT]], %[[C0]] : i16
+// CHECK-NEXT:   %[[LHS_SUB:.*]] = arith.subi %[[C32767]], %[[LHS_INT]] : i16
+// CHECK-NEXT:   %[[LHS_SELECT:.*]] = arith.select %[[LHS_CMP]], %[[LHS_SUB]], %[[LHS_INT]] : i16
+// CHECK-NEXT:   %[[RHS_INT:.*]] = arith.bitcast %[[RHS_IN]] : bf16 to i16
+// CHECK-NEXT:   %[[RHS_CMP:.*]] = arith.cmpi slt, %[[RHS_INT]], %[[C0]] : i16
+// CHECK-NEXT:   %[[RHS_SUB:.*]] = arith.subi %[[C32767]], %[[RHS_INT]] : i16
+// CHECK-NEXT:   %[[RHS_SELECT:.*]] = arith.select %[[RHS_CMP]], %[[RHS_SUB]], %[[RHS_INT]] : i16
+// CHECK-NEXT:   %[[RESULT:.*]] = arith.cmpi slt, %[[LHS_SELECT]], %[[RHS_SELECT]] : i16
+// CHECK-NEXT:   linalg.yield %[[RESULT]] : i1
+
+// CHECK-PRIMITIVE-DAG: %[[C0:.*]] = arith.constant 0 : i16
+// CHECK-PRIMITIVE-DAG: %[[C32767:.*]] = arith.constant 32767 : i16
+// CHECK-PRIMITIVE: linalg.map
+// CHECK-PRIMITIVE-SAME: ins(
+// CHECK-PRIMITIVE-SAME: outs(
+// CHECK-PRIMITIVE-NEXT: (%[[LHS_IN:[a-zA-Z0-9]*]]: bf16, %[[RHS_IN:.*]]: bf16) {
+// CHECK-PRIMITIVE-NEXT:   %[[LHS_INT:.*]] = arith.bitcast %[[LHS_IN]] : bf16 to i16
+// CHECK-PRIMITIVE-NEXT:   %[[LHS_CMP:.*]] = arith.cmpi slt, %[[LHS_INT]], %[[C0]] : i16
+// CHECK-PRIMITIVE-NEXT:   %[[LHS_SUB:.*]] = arith.subi %[[C32767]], %[[LHS_INT]] : i16
+// CHECK-PRIMITIVE-NEXT:   %[[LHS_SELECT:.*]] = arith.select %[[LHS_CMP]], %[[LHS_SUB]], %[[LHS_INT]] : i16
+// CHECK-PRIMITIVE-NEXT:   %[[RHS_INT:.*]] = arith.bitcast %[[RHS_IN]] : bf16 to i16
+// CHECK-PRIMITIVE-NEXT:   %[[RHS_CMP:.*]] = arith.cmpi slt, %[[RHS_INT]], %[[C0]] : i16
+// CHECK-PRIMITIVE-NEXT:   %[[RHS_SUB:.*]] = arith.subi %[[C32767]], %[[RHS_INT]] : i16
+// CHECK-PRIMITIVE-NEXT:   %[[RHS_SELECT:.*]] = arith.select %[[RHS_CMP]], %[[RHS_SUB]], %[[RHS_INT]] : i16
+// CHECK-PRIMITIVE-NEXT:   %[[RESULT:.*]] = arith.cmpi slt, %[[LHS_SELECT]], %[[RHS_SELECT]] : i16
+// CHECK-PRIMITIVE-NEXT:   linalg.yield %[[RESULT]] : i1
+
+// -----
+
+// CHECK-LABEL: func @int_cmp
+// CHECK-PRIMITIVE-LABEL: func @int_cmp
+func.func @int_cmp(%lhs: tensor<2x2xi32>,
+              %rhs: tensor<2x2xi32>) -> tensor<2x2xi1> {
+  %0 = "stablehlo.compare"(%lhs, %rhs) {comparison_direction = #stablehlo<comparison_direction LT>}
+          : (tensor<2x2xi32>, tensor<2x2xi32>) -> (tensor<2x2xi1>)
+  func.return %0 : tensor<2x2xi1>
+}
+// CHECK: tensor.empty() : tensor<2x2xi1>
+// CHECK: linalg.generic
+// CHECK-NEXT: ^bb0(%[[LHS_IN:.*]]: i32, %[[RHS_IN:.*]]: i32, %{{.*}}: i1):
+// CHECK-NEXT:   %[[RESULT:.*]] = arith.cmpi slt, %[[LHS_IN]], %[[RHS_IN]] : i32
+// CHECK-NEXT:   linalg.yield %[[RESULT]] : i1
+// CHECK-PRIMITIVE: linalg.map
+// CHECK-PRIMITIVE: arith.cmpi
+
+// -----
+
+// CHECK-LABEL: func @complex_cmp_eq
+// CHECK-PRIMITIVE-LABEL: func @complex_cmp_eq
+func.func @complex_cmp_eq(%lhs: tensor<2xcomplex<f32>>,
+                     %rhs: tensor<2xcomplex<f32>>) -> tensor<2xi1> {
+  %0 = "stablehlo.compare"(%lhs, %rhs) {comparison_direction = #stablehlo<comparison_direction EQ>}
+          : (tensor<2xcomplex<f32>>, tensor<2xcomplex<f32>>) -> (tensor<2xi1>)
+  func.return %0 : tensor<2xi1>
+}
+// CHECK: tensor.empty() : tensor<2xi1>
+// CHECK: linalg.generic
+// CHECK-NEXT: ^bb0(%[[LHS_IN:.*]]: complex<f32>, %[[RHS_IN:.*]]: complex<f32>, %[[RESULT_OUT:.*]]: i1):
+// CHECK-NEXT:   %[[RESULT:.*]] = complex.eq %[[LHS_IN]], %[[RHS_IN]] : complex<f32>
+// CHECK-NEXT:   linalg.yield %[[RESULT]] : i1
+// CHECK-PRIMITIVE: linalg.map
+// CHECK-PRIMITIVE: complex.eq
+
+// -----
+
+// CHECK-LABEL: func @complex_cmp_neq
+// CHECK-PRIMITIVE-LABEL: func @complex_cmp_neq
+func.func @complex_cmp_neq(%lhs: tensor<2xcomplex<f64>>,
+                      %rhs: tensor<2xcomplex<f64>>) -> tensor<2xi1> {
+  %0 = "stablehlo.compare"(%lhs, %rhs) {comparison_direction = #stablehlo<comparison_direction NE>}
+          : (tensor<2xcomplex<f64>>, tensor<2xcomplex<f64>>) -> (tensor<2xi1>)
+  func.return %0 : tensor<2xi1>
+}
+// CHECK: tensor.empty() : tensor<2xi1>
+// CHECK: linalg.generic
+// CHECK-NEXT: ^bb0(%[[LHS_IN:.*]]: complex<f64>, %[[RHS_IN:.*]]: complex<f64>, %[[RESULT_OUT:.*]]: i1):
+// CHECK-NEXT:   %[[RESULT:.*]] = complex.neq %[[LHS_IN]], %[[RHS_IN]] : complex<f64>
+// CHECK-NEXT:   linalg.yield %[[RESULT]] : i1
+// CHECK-PRIMITIVE: linalg.map
+// CHECK-PRIMITIVE: complex.neq
+
+// -----
+
+// CHECK-LABEL: func @float_cos
+// CHECK-PRIMITIVE-LABEL: func @float_cos
+func.func @float_cos(%arg0: tensor<2x2xf32>) -> tensor<2x2xf32> {
+  // CHECK: linalg.generic
+  // CHECK: math.cos
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: math.cos
+  %0 = "stablehlo.cosine"(%arg0) : (tensor<2x2xf32>) -> tensor<2x2xf32>
+  func.return %0 : tensor<2x2xf32>
+}
+
+// -----
+
+// CHECK-LABEL: func @complex_cos
+// CHECK-PRIMITIVE-LABEL: func @complex_cos
+func.func @complex_cos(%arg0: tensor<2x2xcomplex<f32>>) -> tensor<2x2xcomplex<f32>> {
+  // CHECK: linalg.generic
+  // CHECK: complex.cos
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: complex.cos
+  %0 = "stablehlo.cosine"(%arg0) : (tensor<2x2xcomplex<f32>>) -> tensor<2x2xcomplex<f32>>
+  func.return %0 : tensor<2x2xcomplex<f32>>
+}
+
+// -----
+
+// CHECK-LABEL: func @float_sin
+// CHECK-PRIMITIVE-LABEL: func @float_sin
+func.func @float_sin(%arg0: tensor<2x2xf32>) -> tensor<2x2xf32> {
+  // CHECK: linalg.generic
+  // CHECK: math.sin
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: math.sin
+  %0 = "stablehlo.sine"(%arg0) : (tensor<2x2xf32>) -> tensor<2x2xf32>
+  func.return %0 : tensor<2x2xf32>
+}
+
+// -----
+
+// CHECK-LABEL: func @complex_sin
+// CHECK-PRIMITIVE-LABEL: func @complex_sin
+func.func @complex_sin(%arg0: tensor<2x2xcomplex<f32>>) -> tensor<2x2xcomplex<f32>> {
+  // CHECK: linalg.generic
+  // CHECK: complex.sin
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: complex.sin
+  %0 = "stablehlo.sine"(%arg0) : (tensor<2x2xcomplex<f32>>) -> tensor<2x2xcomplex<f32>>
+  func.return %0 : tensor<2x2xcomplex<f32>>
+}
+
+// -----
+
+// CHECK-LABEL: func @is_finte
+// CHECK-PRIMITIVE-LABEL: func @is_finte
+func.func @is_finte(%input: tensor<2x2xf32>) -> tensor<2x2xi1> {
+  %0 = "stablehlo.is_finite"(%input) : (tensor<2x2xf32>) -> tensor<2x2xi1>
+  func.return %0 : tensor<2x2xi1>
+}
+// CHECK: %[[POS_INF:.+]] = arith.constant 0x7F800000 : f32
+// CHECK: linalg.generic
+// CHECK-NEXT: ^bb0(%[[OPERAND_IN:.*]]: f32
+// CHECK-NEXT:   %[[ABS_X:.+]] = math.absf %[[OPERAND_IN]] : f32
+// CHECK-NEXT:   %[[RESULT:.+]] = arith.cmpf one, %[[ABS_X]], %[[POS_INF]] : f32
+// CHECK-NEXT:   linalg.yield %[[RESULT]] : i1
+
+// CHECK-PRIMITIVE: linalg.map
+// CHECK-PRIMITIVE: math.absf
+// CHECK-PRIMITIVE: arith.cmpf
+
+// -----
+
+// CHECK-LABEL: func @round_nearest_even
+// CHECK-PRIMITIVE-LABEL: func @round_nearest_even
+func.func @round_nearest_even(%val: tensor<2x2xf32>) -> tensor<2x2xf32> {
+  // CHECK: ^{{[a-z0-9_]*}}
+  // CHECK-SAME: %[[IN:[a-zA-Z0-9_]*]]: f32
+  // CHECK-SAME: %[[OUT:[a-zA-Z0-9_]*]]: f32
+  // CHECK: %[[ROUND:.+]] = math.roundeven %[[IN]]
+  // CHECK: linalg.yield %[[ROUND]]
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: math.roundeven
+  %0 = "stablehlo.round_nearest_even"(%val) : (tensor<2x2xf32>) -> (tensor<2x2xf32>)
+  func.return %0 : tensor<2x2xf32>
+}
+
+// -----
+
+// CHECK-LABEL: func @round
+// CHECK-PRIMITIVE-LABEL: func @round
+func.func @round(%val: tensor<2x2xf32>) -> tensor<2x2xf32> {
+  // CHECK: ^{{[a-z0-9_]*}}
+  // CHECK-SAME: %[[IN:[a-zA-Z0-9_]*]]: f32
+  // CHECK-SAME: %[[OUT:[a-zA-Z0-9_]*]]: f32
+  // CHECK: %[[ROUND:.+]] = math.round %[[IN]]
+  // CHECK: linalg.yield %[[ROUND]]
+  // CHECK-PRIMITIVE: linalg.map
+  // CHECK-PRIMITIVE: math.round
+  %0 = "stablehlo.round_nearest_afz"(%val) : (tensor<2x2xf32>) -> (tensor<2x2xf32>)
+  func.return %0 : tensor<2x2xf32>
+}
+
+// -----
+
+// CHECK-LABEL: func @select
+func.func @select(%pred: tensor<2x2xi1>, %lhs: tensor<2x2xf32>,
+             %rhs: tensor<2x2xf32>) -> tensor<2x2xf32> {
+  %0 = "stablehlo.select"(%pred, %lhs, %rhs)
+         : (tensor<2x2xi1>, tensor<2x2xf32>, tensor<2x2xf32>) -> (tensor<2x2xf32>)
+  func.return %0 : tensor<2x2xf32>
+}
+// CHECK: tensor.empty() : tensor<2x2xf32>
+// CHECK: linalg.generic
+// CHECK-NEXT: ^bb0(%[[PRED_IN:.*]]: i1, %[[LHS_IN:.*]]: f32, %[[RHS_IN:.*]]: f32, %{{.*}}: f32):
+// CHECK-NEXT:   %[[RESULT:.*]] = arith.select %[[PRED_IN]], %[[LHS_IN]], %[[RHS_IN]] : f32
+// CHECK-NEXT:   linalg.yield %[[RESULT]] : f32
+
+// CHECK-PRIMITIVE-LABEL: func @select
+// CHECK-PRIMITIVE: tensor.empty() : tensor<2x2xf32>
+// CHECK-PRIMITIVE: linalg.map { arith.select }
+// CHECK-PRIMITIVE-SAME: ins(
+// CHECK-PRIMITIVE-SAME: outs(
+
+// -----
+
+// CHECK-DAG:   #[[SCALAR_MAP:.*]] = affine_map<(d0, d1) -> ()>
+// CHECK-DAG:   #[[ID_MAP:.*]] = affine_map<(d0, d1) -> (d0, d1)>
+// CHECK-LABEL: func @select_scalar_pred_dyn
+// CHECK-SAME:  (%[[PRED:.*]]: tensor<i1>, %[[LHS:.*]]: tensor<2x?xf32>, %[[RHS:.*]]: tensor<2x?xf32>)
+func.func @select_scalar_pred_dyn(%pred : tensor<i1>, %lhs: tensor<2x?xf32>, %rhs: tensor<2x?xf32>) -> tensor<2x?xf32> {
+  %0 = "stablehlo.select"(%pred, %lhs, %rhs) {someattr} : (tensor<i1>, tensor<2x?xf32>, tensor<2x?xf32>) -> (tensor<2x?xf32>)
+  func.return %0 : tensor<2x?xf32>
+}
+// CHECK-DAG:  %[[C1:.*]] = arith.constant 1
+// CHECK-DAG:  %[[DIM:.*]] = tensor.dim %[[LHS]], %[[C1]]
+// CHECK-DAG:  %[[DST:.*]] = tensor.empty(%[[DIM]])
+// CHECK:      linalg.generic
+// CHECK-SAME:   indexing_maps = [#[[SCALAR_MAP]], #[[ID_MAP]], #[[ID_MAP]], #[[ID_MAP]]]
+// CHECK-SAME:   iterator_types = ["parallel", "parallel"]
+// CHECK-SAME:   ins(%[[PRED]], %[[LHS]], %[[RHS]] : tensor<i1>, tensor<2x?xf32>, tensor<2x?xf32>)
+// CHECK-SAME:   outs(%[[DST]] : tensor<2x?xf32>)
+// CHECK-SAME:   {someattr}
+// CHECK:      ^bb0(%[[PRED_:.*]]: i1, %[[LHS_:.*]]: f32, %[[RHS_:.*]]: f32, %{{.*}}: f32):
+// CHECK:        %[[RES:.*]] = arith.select %[[PRED_]], %[[LHS_]], %[[RHS_]] : f32
+// CHECK:        linalg.yield %[[RES]]
+
+// CHECK-PRIMITIVE-LABEL: func @select_scalar_pred_dyn
+// CHECK-PRIMITIVE-SAME:  (%[[PRED:.*]]: tensor<i1>, %[[LHS:.*]]: tensor<2x?xf32>, %[[RHS:.*]]: tensor<2x?xf32>)
+// CHECK-PRIMITIVE-DAG:  %[[C1:.*]] = arith.constant 1
+// CHECK-PRIMITIVE-DAG:  %[[DIM:.*]] = tensor.dim %[[LHS]], %[[C1]]
+// CHECK-PRIMITIVE-DAG:  %[[DST:.*]] = tensor.empty(%[[DIM]])
+// CHECK-PRIMITIVE-DAG:  %[[PRED_ELEM:.*]] = tensor.extract %[[PRED]]
+// CHECK-PRIMITIVE:      linalg.map
+// CHECK-PRIMITIVE-SAME:   ins(%[[LHS]], %[[RHS]] : tensor<2x?xf32>, tensor<2x?xf32>)
+// CHECK-PRIMITIVE-SAME:   outs(%[[DST]] : tensor<2x?xf32>)
+// CHECK-PRIMITIVE-SAME:   {someattr}
+// CHECK-PRIMITIVE:      (%[[LHS_:.*]]: f32, %[[RHS_:.*]]: f32) {
+// CHECK-PRIMITIVE:        %[[RES:.*]] = arith.select %[[PRED_ELEM]], %[[LHS_]], %[[RHS_]] : f32
+// CHECK-PRIMITIVE:        linalg.yield %[[RES]]
+
+// -----
+
+// CHECK-LABEL: func @select_mixed
+func.func @select_mixed(%pred: tensor<2x?xi1>, %lhs: tensor<?x2xf32>,
+             %rhs: tensor<2x2xf32>) -> tensor<?x2xf32> {
+  %0 = "stablehlo.select"(%pred, %lhs, %rhs)
+         : (tensor<2x?xi1>, tensor<?x2xf32>, tensor<2x2xf32>) -> (tensor<?x2xf32>)
+  func.return %0 : tensor<?x2xf32>
+}
+
+// -----
+
+// CHECK-LABEL: func @bitcast_convert
+func.func @bitcast_convert(%input: tensor<2x2xi32>) -> tensor<2x2xf32> {
+  %result = "stablehlo.bitcast_convert"(%input) : (tensor<2x2xi32>) -> tensor<2x2xf32>
+  func.return %result : tensor<2x2xf32>
+}
+// CHECK: tensor.empty
+// CHECK: linalg.generic
+// CHECK-NEXT: ^bb0(%[[OPERAND_IN:.*]]: i32, %{{.*}}: f32):
+// CHECK-NEXT:   %[[RESULT:.*]] = arith.bitcast %[[OPERAND_IN]] : i32 to f32
+// CHECK-NEXT:   linalg.yield %[[RESULT]] : f32
+
+// CHECK-PRIMITIVE: linalg.map
+// CHECK-PRIMITIVE: arith.bitcast
+
+// -----
+
+// CHECK-LABEL: func @bitcast_convert_dynamic
+func.func @bitcast_convert_dynamic(%input: tensor<?x?xi32>) -> tensor<?x?xf32> {
+  %result = "stablehlo.bitcast_convert"(%input) : (tensor<?x?xi32>) -> tensor<?x?xf32>
+  func.return %result : tensor<?x?xf32>
+}
+// CHECK: tensor.empty
+// CHECK: linalg.generic
+// CHECK-NEXT: ^bb0(%[[OPERAND_IN:.*]]: i32, %{{.*}}: f32):
+// CHECK-NEXT:   %[[RESULT:.*]] = arith.bitcast %[[OPERAND_IN]] : i32 to f32
+// CHECK-NEXT:   linalg.yield %[[RESULT]] : f32
+
+// CHECK-PRIMITIVE: linalg.map
+// CHECK-PRIMITIVE: arith.bitcast
+
+// -----
+
+// CHECK: #[[MAP0:.*]] = affine_map<(d0, d1) -> (d0)>
+// CHECK: #[[MAP1:.*]] = affine_map<(d0, d1) -> (d0, d1)>
+// CHECK-LABEL: func @bitcast_convert_expand
+func.func @bitcast_convert_expand(%input: tensor<6xi32>) -> tensor<6x4xi8> {
+  %result = "stablehlo.bitcast_convert"(%input) : (tensor<6xi32>) -> tensor<6x4xi8>
+  func.return %result : tensor<6x4xi8>
+}
+
+// CHECK: %[[C8:.*]] = arith.constant 8 : i32
+// CHECK: tensor.empty() : tensor<6x4xi8>
+// CHECK: %[[RESULT:.*]] = linalg.generic {
+// CHECK:    indexing_maps = [#[[MAP0]], #[[MAP1]]],
+// CHECK:    iterator_types = ["parallel", "parallel"]}
+// CHECK:    ^bb0(%[[IN:.*]]: i32, %[[OUT:.*]]: i8):
+// CHECK:      %[[IOTA:.*]] = linalg.index 1 : index
+// CHECK:      %[[IOTA_CASTED:.*]] = arith.index_cast %[[IOTA]] : index to i32
+// CHECK:      %[[AMT:.*]] = arith.muli %[[IOTA_CASTED]], %[[C8]] : i32
+// CHECK:      %[[SHIFT:.*]] = arith.shrui %[[IN]], %[[AMT]] : i32
+// CHECK:      %[[TRUNC:.*]] = arith.trunci %[[SHIFT]] : i32 to i8
+// CHECK:      linalg.yield %[[TRUNC]] : i8
+// CHECK:    } -> tensor<6x4xi8>
+// CHECK:    return %[[RESULT]] : tensor<6x4xi8>
+
+// -----
+
+// CHECK: #[[MAP0:.*]] = affine_map<(d0, d1) -> (d0, d1)>
+// CHECK: #[[MAP1:.*]] = affine_map<(d0, d1) -> (d0)>
+// CHECK-LABEL: func @bitcast_convert_contract
+func.func @bitcast_convert_contract(%input: tensor<7x4xi8>) -> tensor<7xi32> {
+  %result = "stablehlo.bitcast_convert"(%input) : (tensor<7x4xi8>) -> tensor<7xi32>
+  func.return %result : tensor<7xi32>
+}
+// CHECK-DAG: %[[C0:.*]] = arith.constant 0 : i32
+// CHECK-DAG: %[[C8:.*]] = arith.constant 8 : i32
+// CHECK: %[[EMPTY:.*]] = tensor.empty() : tensor<7xi32>
+// CHECK: linalg.fill ins(%[[C0]] : i32) outs(%[[EMPTY]] : tensor<7xi32>) -> tensor<7xi32>
+// CHECK: %[[RESULT:.*]] = linalg.generic {
+// CHECK:    indexing_maps = [#[[MAP0]], #[[MAP1]]],
+// CHECK:    iterator_types = ["parallel", "reduction"]}
+// CHECK:    ^bb0(%[[IN:.*]]: i8, %[[OUT:.*]]: i32):
+// CHECK:      %[[IOTA:.*]] = linalg.index 1 : index
+// CHECK:      %[[IOTA_CASTED:.*]] = arith.index_cast %[[IOTA]] : index to i32
+// CHECK:      %[[AMT:.*]] = arith.muli %[[IOTA_CASTED]], %[[C8]] : i3
+// CHECK:      %[[EXT:.*]] = arith.extui %[[IN]] : i8 to i32
+// CHECK:      %[[SHIFT:.*]] = arith.shli %[[EXT]], %[[AMT]] : i32
+// CHECK:      %[[OR:.*]] = arith.ori %[[SHIFT]], %[[OUT]] : i32
+// CHECK:      linalg.yield %[[OR]] : i32
+// CHECK: } -> tensor<7xi32>
+// CHECK: return %[[RESULT]] : tensor<7xi32>
+
+// -----
+
+// CHECK-LABEL:   func @concatenate(
+// CHECK-SAME:   %[[VAL_0:[a-zA-Z0-9_]*]]
+// CHECK-SAME:   %[[VAL_1:[a-zA-Z0-9_]*]]
+// CHECK-SAME:   %[[VAL_2:[a-zA-Z0-9_]*]]
+// CHECK-DAG:       %[[C0:.*]] = arith.constant 0 : index
+// CHECK-DAG:       %[[C1:.*]] = arith.constant 1 : index
+// CHECK:           %[[VAL_5:.*]] = tensor.dim %[[VAL_0]], %[[C0]] : tensor<?x?xi32>
+// CHECK:           %[[VAL_7:.*]] = tensor.dim %[[VAL_0]], %[[C1]] : tensor<?x?xi32>
+// CHECK:           %[[VAL_9:.*]] = tensor.dim %[[VAL_1]], %[[C1]] : tensor<?x?xi32>
+// CHECK:           %[[VAL_15:.*]] = tensor.dim %[[VAL_2]], %[[C1]] : tensor<?x?xi32>
+// CHECK:           %[[VAL_16:.*]] = arith.addi %[[VAL_7]], %[[VAL_9]] : index
+// CHECK:           %[[VAL_17:.*]] = arith.addi %[[VAL_16]], %[[VAL_15]] : index
+// CHECK:           %[[VAL_23:.*]] = tensor.empty(%[[VAL_5]], %[[VAL_17]]) : tensor<?x?xi32>
+// CHECK:           %[[VAL_24:.*]] = linalg.generic {indexing_maps = [#map], iterator_types = ["parallel", "parallel"]} outs(%[[VAL_23]] : tensor<?x?xi32>) {
+// CHECK:           ^bb0(%[[VAL_25:.*]]: i32):
+// CHECK:             %[[VAL_26:.*]] = linalg.index 0 : index
+// CHECK:             %[[VAL_28:.*]] = linalg.index 1 : index
+// CHECK:             %[[VAL_30:.*]] = tensor.dim %[[VAL_0]], %[[C1]] : tensor<?x?xi32>
+// CHECK:             %[[VAL_32:.*]] = arith.cmpi ult, %[[VAL_28]], %[[VAL_30]] : index
+// CHECK:             %[[VAL_33:.*]] = scf.if %[[VAL_32]] -> (i32) {
+// CHECK:               %[[VAL_35:.*]] = tensor.extract %[[VAL_0]][%[[VAL_26]], %[[VAL_28]]] : tensor<?x?xi32>
+// CHECK:               scf.yield %[[VAL_35]] : i32
+// CHECK:             } else {
+// CHECK:               %[[VAL_37:.*]] = tensor.dim %[[VAL_1]], %[[C1]] : tensor<?x?xi32>
+// CHECK:               %[[VAL_38:.*]] = arith.addi %[[VAL_30]], %[[VAL_37]] : index
+// CHECK:               %[[VAL_39:.*]] = arith.cmpi ult, %[[VAL_28]], %[[VAL_38]] : index
+// CHECK:               %[[VAL_40:.*]] = scf.if %[[VAL_39]] -> (i32) {
+// CHECK:                 %[[VAL_41:.*]] = arith.subi %[[VAL_28]], %[[VAL_30]] : index
+// CHECK:                 %[[VAL_42:.*]] = tensor.extract %[[VAL_1]][%[[VAL_26]], %[[VAL_41]]] : tensor<?x?xi32>
+// CHECK:                 scf.yield %[[VAL_42]] : i32
+// CHECK:               } else {
+// CHECK:                 %[[VAL_43:.*]] = arith.subi %[[VAL_28]], %[[VAL_38]] : index
+// CHECK:                 %[[VAL_44:.*]] = tensor.extract %[[VAL_2]][%[[VAL_26]], %[[VAL_43]]] : tensor<?x?xi32>
+// CHECK:                 scf.yield %[[VAL_44]] : i32
+// CHECK:               }
+// CHECK:               scf.yield %[[VAL_45:.*]] : i32
+// CHECK:             }
+// CHECK:             linalg.yield %[[VAL_46:.*]] : i32
+// CHECK:           } -> tensor<?x?xi32>
+// CHECK:           return %[[VAL_47:.*]] : tensor<?x?xi32>
+// CHECK:         }
+func.func @concatenate(%a: tensor<?x?xi32>, %b: tensor<?x?xi32>, %c: tensor<?x?xi32>) -> tensor<?x?xi32> {
+    %concat = "stablehlo.concatenate"(%a, %b, %c) {
+      dimension = 1
+    } : (tensor<?x?xi32>, tensor<?x?xi32>, tensor<?x?xi32>) -> tensor<?x?xi32>
+    func.return %concat : tensor<?x?xi32>
+}
+
+// -----
+
+// CHECK-LABEL:   func @concatenate_unsigned(
+// CHECK-SAME:   %[[VAL_0:[a-zA-Z0-9_]*]]
+// CHECK-SAME:   %[[VAL_1:[a-zA-Z0-9_]*]]
+// CHECK-SAME:   %[[VAL_2:[a-zA-Z0-9_]*]]
+// CHECK-DAG:       %[[C0:.*]] = arith.constant 0 : index
+// CHECK-DAG:       %[[C1:.*]] = arith.constant 1 : index
+// CHECK-DAG:       %[[VAL_5:.*]] = builtin.unrealized_conversion_cast %[[VAL_2]] : tensor<?x?xui32> to tensor<?x?xi32>
+// CHECK-DAG:       %[[VAL_4:.*]] = builtin.unrealized_conversion_cast %[[VAL_1]] : tensor<?x?xui32> to tensor<?x?xi32>
+// CHECK-DAG:       %[[VAL_3:.*]] = builtin.unrealized_conversion_cast %[[VAL_0]] : tensor<?x?xui32> to tensor<?x?xi32>
+// CHECK:           %[[VAL_8:.*]] = tensor.dim %[[VAL_3]], %[[C0]] : tensor<?x?xi32>
+// CHECK:           %[[VAL_10:.*]] = tensor.dim %[[VAL_3]], %[[C1]] : tensor<?x?xi32>
+// CHECK:           %[[VAL_14:.*]] = tensor.dim %[[VAL_4]], %[[C1]] : tensor<?x?xi32>
+// CHECK:           %[[VAL_18:.*]] = tensor.dim %[[VAL_5]], %[[C1]] : tensor<?x?xi32>
+// CHECK:           %[[VAL_19:.*]] = arith.addi %[[VAL_10]], %[[VAL_14]] : index
+// CHECK:           %[[VAL_20:.*]] = arith.addi %[[VAL_19]], %[[VAL_18]] : index
+// CHECK:           %[[VAL_26:.*]] = tensor.empty(%[[VAL_8]], %[[VAL_20]]) : tensor<?x?xi32>
+// CHECK:           %[[VAL_27:.*]] = linalg.generic {indexing_maps = [#map], iterator_types = ["parallel", "parallel"]} outs(%[[VAL_26]] : tensor<?x?xi32>) {
+// CHECK:           ^bb0(%[[VAL_28:.*]]: i32):
+// CHECK:             %[[VAL_29:.*]] = linalg.index 0 : index
+// CHECK:             %[[VAL_30:.*]] = linalg.index 1 : index
+// CHECK:             %[[VAL_33:.*]] = tensor.dim %[[VAL_3]], %[[C1]] : tensor<?x?xi32>
+// CHECK:             %[[VAL_35:.*]] = arith.cmpi ult, %[[VAL_30]], %[[VAL_33]] : index
+// CHECK:             %[[VAL_36:.*]] = scf.if %[[VAL_35]] -> (i32) {
+// CHECK:               %[[VAL_38:.*]] = tensor.extract %[[VAL_3]][%[[VAL_29]], %[[VAL_30]]] : tensor<?x?xi32>
+// CHECK:               scf.yield %[[VAL_38]] : i32
+// CHECK:             } else {
+// CHECK:               %[[VAL_40:.*]] = tensor.dim %[[VAL_4]], %[[C1]] : tensor<?x?xi32>
+// CHECK:               %[[VAL_41:.*]] = arith.addi %[[VAL_33]], %[[VAL_40]] : index
+// CHECK:               %[[VAL_42:.*]] = arith.cmpi ult, %[[VAL_30]], %[[VAL_41]] : index
+// CHECK:               %[[VAL_43:.*]] = scf.if %[[VAL_42]] -> (i32) {
+// CHECK:                 %[[VAL_44:.*]] = arith.subi %[[VAL_30]], %[[VAL_33]] : index
+// CHECK:                 %[[VAL_45:.*]] = tensor.extract %[[VAL_4]][%[[VAL_29]], %[[VAL_44]]] : tensor<?x?xi32>
+// CHECK:                 scf.yield %[[VAL_45]] : i32
+// CHECK:               } else {
+// CHECK:                 %[[VAL_46:.*]] = arith.subi %[[VAL_30]], %[[VAL_41]] : index
+// CHECK:                 %[[VAL_47:.*]] = tensor.extract %[[VAL_5]][%[[VAL_29]], %[[VAL_46]]] : tensor<?x?xi32>
+// CHECK:                 scf.yield %[[VAL_47]] : i32
+// CHECK:               }
+// CHECK:               scf.yield %[[VAL_48:.*]] : i32
+// CHECK:             }
+// CHECK:             linalg.yield %[[VAL_49:.*]] : i32
+// CHECK:           } -> tensor<?x?xi32>
+// CHECK:           %[[VAL_50:.*]] = builtin.unrealized_conversion_cast %[[VAL_51:.*]] : tensor<?x?xi32> to tensor<?x?xui32>
+// CHECK:           return %[[VAL_50]] : tensor<?x?xui32>
+// CHECK:         }
+func.func @concatenate_unsigned(%a: tensor<?x?xui32>, %b: tensor<?x?xui32>, %c: tensor<?x?xui32>) -> tensor<?x?xui32> {
+    %concat = "stablehlo.concatenate"(%a, %b, %c) {
+      dimension = 1
+    } : (tensor<?x?xui32>, tensor<?x?xui32>, tensor<?x?xui32>) -> tensor<?x?xui32>
+    func.return %concat : tensor<?x?xui32>
+}
+
+// -----
+
+// CHECK-DAG: #[[MAP0:.+]] = affine_map<(d0, d1, d2, d3) -> (d0, d1, d2)>
+// CHECK-DAG: #[[MAP1:.+]] = affine_map<(d0, d1, d2, d3) -> (d0, d2, d3)>
+// CHECK-DAG: #[[MAP2:.+]] = affine_map<(d0, d1, d2, d3) -> (d0, d1, d3)>
+// CHECK-LABEL: func @einsum_basic
+func.func @einsum_basic(%arg0: tensor<3x4x5xf32>, %arg1: tensor<3x5x6xf32>) -> tensor<3x4x6xf32> {
+  %0 = "stablehlo.einsum"(%arg0, %arg1) {einsum_config = "ijk,ikm->ijm", someattr}: (tensor<3x4x5xf32>, tensor<3x5x6xf32>) -> tensor<3x4x6xf32>
+  func.return %0 : tensor<3x4x6xf32>
+}
+// CHECK-SAME:  (%[[LHS:.*]]: tensor<3x4x5xf32>, %[[RHS:.*]]: tensor<3x5x6xf32>)
+// CHECK-DAG: %[[INIT:.*]] = tensor.empty() : tensor<3x4x6xf32>
+// CHECK-DAG: %[[ZERO:.*]] = arith.constant 0.000000e+00 : f32
+// CHECK: %[[FILL:.*]] = linalg.fill ins(%[[ZERO]]{{.*}}outs(%[[INIT]]
+// CHECK: linalg.generic
+// CHECK-SAME: indexing_maps = [#[[MAP0]], #[[MAP1]], #[[MAP2]]]
+// CHECK-SAME: iterator_types = ["parallel", "parallel", "reduction", "parallel"]
+// CHECK-SAME: ins(%[[LHS]], %[[RHS]] : tensor<3x4x5xf32>, tensor<3x5x6xf32>)
+// CHECK-SAME: outs(%[[FILL]] : tensor<3x4x6xf32>)
+// CHECK-SAME: {someattr}
+// CHECK: ^bb0(%[[LHS_:.*]]: f32, %[[RHS_:.*]]: f32, %[[OUT_:.*]]: f32):
+// CHECK:   %[[MUL:.*]] = arith.mulf %[[LHS_]], %[[RHS_]] : f32
+// CHECK:   %[[RES:.*]] = arith.addf %[[OUT_]], %[[MUL]] : f32
+// CHECK:   linalg.yield %[[RES]]
+
+// -----
+
+// CHECK: #[[MAP0:.+]] = affine_map<(d0, d1, d2) -> (d0, d1, d2)>
+// CHECK: func @einsum_pointwisemul
+func.func @einsum_pointwisemul(%arg0: tensor<3x4x5xf32>, %arg1: tensor<3x4x5xf32>) -> tensor<3x4x5xf32> {
+  %0 = "stablehlo.einsum"(%arg0, %arg1) {einsum_config = "abc,abc->abc"} : (tensor<3x4x5xf32>, tensor<3x4x5xf32>) -> tensor<3x4x5xf32>
+  func.return %0 : tensor<3x4x5xf32>
+}
+// CHECK-SAME:  (%[[LHS:.*]]: tensor<3x4x5xf32>, %[[RHS:.*]]: tensor<3x4x5xf32>)
+// CHECK: tensor.empty() : tensor<3x4x5xf32>
+// CHECK: linalg.generic
+// CHECK-SAME: indexing_maps = [#[[MAP0]], #[[MAP0]], #[[MAP0]]]
+// CHECK-SAME: iterator_types = ["parallel", "parallel", "parallel"]
+// CHECK-SAME: ins(%[[LHS]], %[[RHS]] : tensor<3x4x5xf32>, tensor<3x4x5xf32>)
+// CHECK-SAME: outs(%[[DST:.+]] : tensor<3x4x5xf32>)
+// CHECK: ^bb0(%[[LHS_:.*]]: f32, %[[RHS_:.*]]: f32, %[[OUT_:.*]]: f32):
+// CHECK:   %[[RES:.*]] = arith.mulf %[[LHS_]], %[[RHS_]] : f32
+// CHECK:   linalg.yield %[[RES]]
+
+// -----
+
+// CHECK-DAG: #[[MAP0:.+]] = affine_map<(d0, d1, d2) -> (d0, d1)>
+// CHECK-DAG: #[[MAP1:.+]] = affine_map<(d0, d1, d2) -> (d1, d2)>
+// CHECK-DAG: #[[MAP2:.+]] = affine_map<(d0, d1, d2) -> (d0, d2)>
+// CHECK: func @einsum_matmul
+func.func @einsum_matmul(%arg0: tensor<7x9xf32>, %arg1: tensor<9x5xf32>) -> tensor<7x5xf32> {
+  %0 = "stablehlo.einsum"(%arg0, %arg1) {einsum_config = "ae,ed->ad"}: (tensor<7x9xf32>, tensor<9x5xf32>) -> tensor<7x5xf32>
+  func.return %0 : tensor<7x5xf32>
+}
+// CHECK-SAME:  (%[[LHS:.*]]: tensor<7x9xf32>, %[[RHS:.*]]: tensor<9x5xf32>)
+// CHECK-DAG: %[[INIT:.*]] = tensor.empty() : tensor<7x5xf32>
+// CHECK-DAG: %[[ZERO:.*]] = arith.constant 0.000000e+00 : f32
+// CHECK: %[[FILL:.*]] = linalg.fill ins(%[[ZERO]]{{.*}}outs(%[[INIT]]
+// CHECK: linalg.generic
+// CHECK-SAME: indexing_maps = [#[[MAP0]], #[[MAP1]], #[[MAP2]]]
+// CHECK-SAME: iterator_types = ["parallel", "reduction", "parallel"]
+// CHECK-SAME: ins(%[[LHS]], %[[RHS]] : tensor<7x9xf32>, tensor<9x5xf32>)
+// CHECK-SAME: outs(%[[FILL]] : tensor<7x5xf32>)
+// CHECK: ^bb0(%[[LHS_:.*]]: f32, %[[RHS_:.*]]: f32, %[[OUT_:.*]]: f32):
+// CHECK:   %[[MUL:.*]] = arith.mulf %[[LHS_]], %[[RHS_]] : f32
+// CHECK:   %[[RES:.*]] = arith.addf %[[OUT_]], %[[MUL]] : f32
+// CHECK:   linalg.yield %[[RES]]
+
+// -----
+
+// CHECK-DAG: #[[MAP0:.+]] = affine_map<(d0, d1, d2, d3, d4, d5) -> (d0, d1, d2, d3, d4)>
+// CHECK-DAG: #[[MAP1:.+]] = affine_map<(d0, d1, d2, d3, d4, d5) -> (d4, d5)>
+// CHECK-DAG: #[[MAP2:.+]] = affine_map<(d0, d1, d2, d3, d4, d5) -> (d0, d1, d2, d3, d5)>
+// CHECK: func @einsum_broadcast4
+func.func @einsum_broadcast4(%arg0: tensor<3x4x5x6x7xf32>, %arg1: tensor<7x8xf32>) -> tensor<3x4x5x6x8xf32> {
+  %0 = "stablehlo.einsum"(%arg0, %arg1) {einsum_config = "abcdh,hg->abcdg"}: (tensor<3x4x5x6x7xf32>, tensor<7x8xf32>) -> tensor<3x4x5x6x8xf32>
+  func.return %0 : tensor<3x4x5x6x8xf32>
+}
+// CHECK-SAME:  (%[[LHS:.*]]: tensor<3x4x5x6x7xf32>, %[[RHS:.*]]: tensor<7x8xf32>)
+// CHECK-DAG: %[[INIT:.*]] = tensor.empty() : tensor<3x4x5x6x8xf32>
+// CHECK-DAG: %[[ZERO:.*]] = arith.constant 0.000000e+00 : f32
+// CHECK: %[[FILL:.*]] = linalg.fill ins(%[[ZERO]]{{.*}}outs(%[[INIT]]
+// CHECK: linalg.generic
+// CHECK-SAME: indexing_maps = [#[[MAP0]], #[[MAP1]], #[[MAP2]]]
+// CHECK-SAME: iterator_types = ["parallel", "parallel", "parallel", "parallel", "reduction", "parallel"]
+// CHECK-SAME: ins(%[[LHS]], %[[RHS]] : tensor<3x4x5x6x7xf32>, tensor<7x8xf32>)
+// CHECK-SAME: outs(%[[FILL]] : tensor<3x4x5x6x8xf32>)
+// CHECK: ^bb0(%[[LHS_:.*]]: f32, %[[RHS_:.*]]: f32, %[[OUT_:.*]]: f32):
+// CHECK:   %[[MUL:.*]] = arith.mulf %[[LHS_]], %[[RHS_]] : f32
+// CHECK:   %[[RES:.*]] = arith.addf %[[OUT_]], %[[MUL]] : f32
+// CHECK:   linalg.yield %[[RES]]
+
+// -----
+
+// CHECK-DAG: #[[MAP0:.+]] = affine_map<(d0, d1, d2, d3) -> (d0, d1, d2)>
+// CHECK-DAG: #[[MAP1:.+]] = affine_map<(d0, d1, d2, d3) -> (d2, d3)>
+// CHECK-DAG: #[[MAP2:.+]] = affine_map<(d0, d1, d2, d3) -> (d0, d1, d3)>
+// CHECK: func @einsum_ellipsis
+func.func @einsum_ellipsis(%arg0: tensor<1x512x128xf32>, %arg1: tensor<128x256xf32>) -> tensor<1x512x256xf32> {
+  %0 = "stablehlo.einsum"(%arg0, %arg1) {einsum_config = "...x,xy->...y"} : (tensor<1x512x128xf32>, tensor<128x256xf32>) -> tensor<1x512x256xf32>
+  func.return %0 : tensor<1x512x256xf32>
+}
+// CHECK-SAME:  (%[[LHS:.*]]: tensor<1x512x128xf32>, %[[RHS:.*]]: tensor<128x256xf32>)
+// CHECK-DAG: %[[INIT:.*]] = tensor.empty() : tensor<1x512x256xf32>
+// CHECK-DAG: %[[ZERO:.*]] = arith.constant 0.000000e+00 : f32
+// CHECK: %[[FILL:.*]] = linalg.fill ins(%[[ZERO]]{{.*}}outs(%[[INIT]]
+// CHECK: linalg.generic
+// CHECK-SAME: indexing_maps = [#[[MAP0]], #[[MAP1]], #[[MAP2]]]
+// CHECK-SAME: iterator_types = ["parallel", "parallel", "reduction", "parallel"]
+// CHECK-SAME: ins(%[[LHS]], %[[RHS]] : tensor<1x512x128xf32>, tensor<128x256xf32>)
+// CHECK-SAME: outs(%[[FILL]] : tensor<1x512x256xf32>)
+// CHECK: ^bb0(%[[LHS_:.*]]: f32, %[[RHS_:.*]]: f32, %[[OUT_:.*]]: f32):
+// CHECK:   %[[MUL:.*]] = arith.mulf %[[LHS_]], %[[RHS_]] : f32
+// CHECK:   %[[RES:.*]] = arith.addf %[[OUT_]], %[[MUL]] : f32
+// CHECK:   linalg.yield %[[RES]]
+
+// -----
+
+// CHECK-DAG: #[[MAP0:.+]] = affine_map<(d0, d1, d2, d3) -> (d0, d1, d2)>
+// CHECK-DAG: #[[MAP1:.+]] = affine_map<(d0, d1, d2, d3) -> (d2, d3)>
+// CHECK-DAG: #[[MAP2:.+]] = affine_map<(d0, d1, d2, d3) -> (d0, d1, d3)>
+// CHECK: func @einsum_dynamic_size_broadcast_dot
+func.func @einsum_dynamic_size_broadcast_dot(%arg0: tensor<?x?x4xf32>, %arg1: tensor<4x?xf32>) -> tensor<?x?x?xf32> {
+  %0 = "stablehlo.einsum"(%arg0, %arg1) {einsum_config = "abc,cd->abd"} : (tensor<?x?x4xf32>, tensor<4x?xf32>) -> tensor<?x?x?xf32>
+  func.return %0 : tensor<?x?x?xf32>
+}
+// CHECK-SAME:  (%[[LHS:.*]]: tensor<?x?x4xf32>, %[[RHS:.*]]: tensor<4x?xf32>)
+// CHECK-DAG: %[[C0:.+]] = arith.constant 0 : index
+// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index
+// CHECK-DAG: %[[ZERO:.*]] = arith.constant 0.000000e+00 : f32
+// CHECK: %[[DIM0:.+]] = tensor.dim %[[LHS]], %[[C0]] : tensor<?x?x4xf32>
+// CHECK: %[[DIM1:.+]] = tensor.dim %[[LHS]], %[[C1]] : tensor<?x?x4xf32>
+// CHECK: %[[DIM2:.+]] = tensor.dim %[[RHS]], %[[C1:.+]] : tensor<4x?xf32>
+// CHECK: %[[INIT:.*]] = tensor.empty(%[[DIM0]], %[[DIM1]], %[[DIM2]]) : tensor<?x?x?xf32>
+// CHECK: %[[FILL:.*]] = linalg.fill ins(%[[ZERO]]{{.*}}outs(%[[INIT]]
+// CHECK: linalg.generic
+// CHECK-SAME: indexing_maps = [#[[MAP0]], #[[MAP1]], #[[MAP2]]]
+// CHECK-SAME: iterator_types = ["parallel", "parallel", "reduction", "parallel"]
+// CHECK-SAME: ins(%[[LHS]], %[[RHS]] : tensor<?x?x4xf32>, tensor<4x?xf32>)
+// CHECK-SAME: outs(%[[FILL]] : tensor<?x?x?xf32>)
+// CHECK: ^bb0(%[[LHS_:.*]]: f32, %[[RHS_:.*]]: f32, %[[OUT_:.*]]: f32):
+// CHECK:   %[[MUL:.*]] = arith.mulf %[[LHS_]], %[[RHS_]] : f32
+// CHECK:   %[[RES:.*]] = arith.addf %[[OUT_]], %[[MUL]] : f32
+// CHECK:   linalg.yield %[[RES]]
+
+// -----
+
+// CHECK-DAG: #[[OPERAND_MAP:.*]] = affine_map<(d0, d1, d2, d3, d4) -> (d4, d0, 0)>
+// CHECK-DAG: #[[RESULT_MAP:.*]] = affine_map<(d0, d1, d2, d3, d4) -> (d0, d1, d2, d3, d4)>
+// CHECK-LABEL: func @broadcast_in_dim
+func.func @broadcast_in_dim(%operand: tensor<5x7x1xf32>) -> tensor<7x10x6x4x5xf32> {
+  %0 = "stablehlo.broadcast_in_dim"(%operand)
+         {broadcast_dimensions = dense<[4,0,2]> : tensor<3xi64>}
+         : (tensor<5x7x1xf32>) -> tensor<7x10x6x4x5xf32>
+  func.return %0 : tensor<7x10x6x4x5xf32>
+}
+// CHECK: tensor.empty() : tensor<7x10x6x4x5xf32>
+// CHECK: linalg.generic {{{.*}}indexing_maps = [#[[OPERAND_MAP]], #[[RESULT_MAP]]]
+// CHECK-NEXT: ^bb0(%[[OPERAND:.*]]: f32, %{{.*}}: f32):
+// CHECK-NEXT:   linalg.yield %[[OPERAND]] : f32
+
+// CHECK-PRIMITIVE-LABEL: func @broadcast_in_dim
+// CHECK-PRIMITIVE: tensor.collapse_shape
+// CHECK-PRIMITIVE: linalg.transpose
+// CHECK-PRIMITIVE:   permutation = [1, 0]
+// CHECK-PRIMITIVE: tensor.empty() : tensor<7x10x6x4x5xf32>
+// CHECK-PRIMITIVE: linalg.broadcast
+// CHECK-PRIMITIVE:   dimensions = [1, 2, 3]
+
+// -----
+
+// CHECK-DAG: #[[OPERAND_MAP:.*]] = affine_map<(d0, d1, d2, d3, d4) -> (d4, d0, 0)>
+// CHECK-DAG: #[[RESULT_MAP:.*]] = affine_map<(d0, d1, d2, d3, d4) -> (d0, d1, d2, d3, d4)>
+// CHECK-LABEL: func @broadcast_in_dim_ui32
+func.func @broadcast_in_dim_ui32(%operand: tensor<5x7x1xui32>) -> tensor<7x10x6x4x5xui32> {
+  %0 = "stablehlo.broadcast_in_dim"(%operand)
+         {broadcast_dimensions = dense<[4,0,2]> : tensor<3xi64>}
+         : (tensor<5x7x1xui32>) -> tensor<7x10x6x4x5xui32>
+  func.return %0 : tensor<7x10x6x4x5xui32>
+}
+// CHECK: builtin.unrealized_conversion_cast %{{.*}} : tensor<5x7x1xui32> to tensor<5x7x1xi32>
+// CHECK: tensor.empty() : tensor<7x10x6x4x5xi32>
+// CHECK: %[[RES:.*]] = linalg.generic {{{.*}}indexing_maps = [#[[OPERAND_MAP]], #[[RESULT_MAP]]]
+// CHECK-NEXT: ^bb0(%[[OPERAND:.*]]: i32, %{{.*}}: i32):
+// CHECK-NEXT:   linalg.yield %[[OPERAND]] : i32
+// CHECK: builtin.unrealized_conversion_cast %[[RES]] : tensor<7x10x6x4x5xi32> to tensor<7x10x6x4x5xui32>
+
+// CHECK-PRIMITIVE-LABEL: func @broadcast_in_dim_ui32
+// CHECK-PRIMITIVE: tensor.collapse_shape
+// CHECK-PRIMITIVE: linalg.transpose
+// CHECK-PRIMITIVE:   permutation = [1, 0]
+// CHECK-PRIMITIVE: tensor.empty() : tensor<7x10x6x4x5xi32>
+// CHECK-PRIMITIVE: %[[RES:.*]] = linalg.broadcast
+// CHECK-PRIMITIVE:   dimensions = [1, 2, 3]
+// CHECK-PRIMITIVE: builtin.unrealized_conversion_cast %[[RES]] : tensor<7x10x6x4x5xi32> to tensor<7x10x6x4x5xui32>
+
+// -----
+
+// CHECK-DAG: #[[OPERAND_MAP:.+]] = affine_map<(d0, d1) -> (d0)>
+// CHECK-DAG: #[[RESULT_MAP:.+]] = affine_map<(d0, d1) -> (d0, d1)>
+// CHECK-LABEL: func @broadcast_in_dim_with_one_to_one
+func.func @broadcast_in_dim_with_one_to_one(
+         %operand: tensor<1xf32>) -> tensor<1x5xf32> {
+  %0 = "stablehlo.broadcast_in_dim"(%operand)
+         {broadcast_dimensions = dense<[0]> : tensor<1xi64>}
+         : (tensor<1xf32>) -> tensor<1x5xf32>
+  func.return %0 : tensor<1x5xf32>
+}
+// CHECK: tensor.empty() : tensor<1x5xf32>
+// CHECK: linalg.generic {{{.*}}indexing_maps = [#[[OPERAND_MAP]], #[[RESULT_MAP]]]
+// CHECK-NEXT: ^bb0(%[[OPERAND:.*]]: f32, %{{.*}}: f32):
+// CHECK-NEXT:   linalg.yield %[[OPERAND]] : f32
+
+// CHECK-PRIMITIVE-LABEL: func @broadcast_in_dim_with_one_to_one
+// CHECK-PRIMITIVE-NOT: tensor.collapse_shape
+// CHECK-PRIMITIVE-NOT: linalg.transpose
+// CHECK-PRIMITIVE:     linalg.broadcast
+// CHECK-PRIMITIVE:       dimensions = [1]
+
+// -----
+
+// CHECK-DAG: #[[OPERAND_MAP:.+]] = affine_map<(d0, d1, d2, d3) -> (d2, d0, d1)>
+// CHECK-DAG: #[[RESULT_MAP:.+]] = affine_map<(d0, d1, d2, d3) -> (d0, d1, d2, d3)>
+// CHECK-LABEL: func @broadcast_in_dim_with_transpose
+func.func @broadcast_in_dim_with_transpose(
+         %operand: tensor<2x3x4xf32>) -> tensor<3x4x2x5xf32> {
+  %0 = "stablehlo.broadcast_in_dim"(%operand)
+         {broadcast_dimensions = dense<[2, 0, 1]> : tensor<3xi64>}
+         : (tensor<2x3x4xf32>) -> tensor<3x4x2x5xf32>
+  func.return %0 : tensor<3x4x2x5xf32>
+}
+// CHECK: tensor.empty() : tensor<3x4x2x5xf32>
+// CHECK: linalg.generic {{{.*}}indexing_maps = [#[[OPERAND_MAP]], #[[RESULT_MAP]]]
+// CHECK-NEXT: ^bb0(%[[OPERAND:.*]]: f32, %{{.*}}: f32):
+// CHECK-NEXT:   linalg.yield %[[OPERAND]] : f32
+
+// CHECK-PRIMITIVE-LABEL: func @broadcast_in_dim_with_transpose
+// CHECK-PRIMITIVE: tensor.empty() : tensor<3x4x2xf32>
+// CHECK-PRIMITIVE: linalg.transpose
+// CHECK-PRIMITIVE:   permutation = [1, 2, 0]
+// CHECK-PRIMITIVE: tensor.empty() : tensor<3x4x2x5xf32>
+// CHECK-PRIMITIVE: linalg.broadcast
+// CHECK-PRIMITIVE:   dimensions = [3]
+
+// -----
+
+// CHECK-DAG: #[[OPERAND_MAP:.*]] = affine_map<(d0, d1, d2) -> ()>
+// CHECK-DAG: #[[RESULT_MAP:.*]] = affine_map<(d0, d1, d2) -> (d0, d1, d2)>
+// CHECK-LABEL: func @broadcast_in_dim_scalar
+func.func @broadcast_in_dim_scalar(%operand: tensor<f32>) -> tensor<7x10x6xf32> {
+  %0 = "stablehlo.broadcast_in_dim"(%operand)
+        {broadcast_dimensions = dense<[]> : tensor<0xi64>}
+        : (tensor<f32>) -> tensor<7x10x6xf32>
+  func.return %0 : tensor<7x10x6xf32>
+}
+// CHECK: tensor.empty() : tensor<7x10x6xf32>
+// CHECK: linalg.generic {{{.*}}indexing_maps = [#[[OPERAND_MAP]], #[[RESULT_MAP]]]
+// CHECK-NEXT: ^bb0(%[[OPERAND:.*]]: f32, %{{.*}}: f32):
+// CHECK-NEXT:   linalg.yield %[[OPERAND]] : f32
+
+// CHECK-PRIMITIVE-LABEL: func @broadcast_in_dim_scalar
+// CHECK-PRIMITIVE: tensor.empty() : tensor<7x10x6xf32>
+// CHECK-PRIMITIVE: linalg.broadcast
+// CHECK-PRIMITIVE:   dimensions = [0, 1, 2]
+
+// -----
+
+// CHECK-DAG: #[[OPERAND_MAP:.+]] = affine_map<(d0, d1, d2) -> ()>
+// CHECK-DAG: #[[RESULT_MAP:.+]] = affine_map<(d0, d1, d2) -> (d0, d1, d2)>
+// CHECK-LABEL: func @broadcast_scalar
+func.func @broadcast_scalar(%arg: tensor<f32>) -> tensor<4x2x1xf32> {
+  %0 = "stablehlo.broadcast"(%arg) {broadcast_sizes = dense<[4, 2, 1]> : tensor<3xi64>} : (tensor<f32>) -> tensor<4x2x1xf32>
+  func.return %0: tensor<4x2x1xf32>
+}
+// CHECK: tensor.empty() : tensor<4x2x1xf32>
+// CHECK: linalg.generic {{{.*}}indexing_maps = [#[[OPERAND_MAP]], #[[RESULT_MAP]]]
+// CHECK-NEXT: ^bb0(%[[OPERAND:.*]]: f32, %{{.*}}: f32):
+// CHECK-NEXT:   linalg.yield %[[OPERAND]] : f32
+
+// CHECK-PRIMITIVE-LABEL: func @broadcast_scalar
+// CHECK-PRIMITIVE: tensor.empty() : tensor<4x2x1xf32>
+// CHECK-PRIMITIVE: linalg.broadcast
+// CHECK-PRIMITIVE-SAME: ins(
+// CHECK-PRIMITIVE-SAME: outs(
+// CHECK-PRIMITIVE-SAME: dimensions = [0, 1, 2]
+
+// -----
+
+// CHECK-DAG: #[[OPERAND_MAP:.+]] = affine_map<(d0, d1, d2, d3, d4, d5) -> (d3, d4, d5)>
+// CHECK-DAG: #[[RESULT_MAP:.+]] = affine_map<(d0, d1, d2, d3, d4, d5) -> (d0, d1, d2, d3, d4, d5)>
+// CHECK-LABEL: func @broadcast
+func.func @broadcast(%arg: tensor<4x?x16xf32>) -> tensor<4x2x1x4x?x16xf32> {
+  %0 = "stablehlo.broadcast"(%arg) {broadcast_sizes = dense<[4, 2, 1]> : tensor<3xi64>} : (tensor<4x?x16xf32>) -> tensor<4x2x1x4x?x16xf32>
+  func.return %0: tensor<4x2x1x4x?x16xf32>
+}
+// CHECK-DAG: %[[C1:.*]] = arith.constant 1 : index
+// CHECK: %[[DIM:.*]] = tensor.dim %{{.*}}, %[[C1]] : tensor<4x?x16xf32>
+// CHECK: %{{.*}} = tensor.empty(%[[DIM]]) : tensor<4x2x1x4x?x16xf32>
+// CHECK: linalg.generic {{{.*}}indexing_maps = [#[[OPERAND_MAP]], #[[RESULT_MAP]]]
+// CHECK-NEXT: ^bb0(%[[OPERAND:.*]]: f32, %{{.*}}: f32):
+// CHECK-NEXT:   linalg.yield %[[OPERAND]] : f32
+
+// CHECK-PRIMITIVE-LABEL: func @broadcast
+// CHECK-PRIMITIVE-DAG: %[[C1:.*]] = arith.constant 1 : index
+// CHECK-PRIMITIVE: %[[DIM:.*]] = tensor.dim %{{.*}}, %[[C1]] : tensor<4x?x16xf32>
+// CHECK-PRIMITIVE: %{{.*}} = tensor.empty(%[[DIM]]) : tensor<4x2x1x4x?x16xf32>
+// CHECK-PRIMITIVE: linalg.broadcast
+// CHECK-PRIMITIVE:   dimensions = [0, 1, 2]
+
+// -----
+
+// CHECK: #[[RESULT_MAP:.*]] = affine_map<(d0, d1) -> (d0, d1)>
+// CHECK-LABEL: func @iota_f32
+func.func @iota_f32() -> tensor<7x10xf32> {
+  %result = "stablehlo.iota"() {iota_dimension = 1 : i64, someattr} : () -> (tensor<7x10xf32>)
+  func.return %result : tensor<7x10xf32>
+}
+// CHECK: tensor.empty
+// CHECK: linalg.generic
+// CHECK-SAME: indexing_maps = [#[[RESULT_MAP]]]
+// CHECK-SAME: {someattr}
+// CHECK-NEXT: ^bb0(%{{.*}}: f32):
+// CHECK-NEXT:   %[[INDEX:.*]] = linalg.index 1
+// CHECK-NEXT:   %[[INT_CAST:.*]] = arith.index_cast %[[INDEX]] : index to i32
+// CHECK-NEXT:   %[[FLOAT_CAST:.*]] = arith.sitofp %[[INT_CAST]] : i32 to f32
+// CHECK-NEXT:   linalg.yield %[[FLOAT_CAST]] : f32
+
+// CHECK-PRIMITIVE-LABEL: func @iota_f32
+// CHECK-PRIMITIVE: %[[EMPTY:.*]] = tensor.empty()
+// CHECK-PRIMITIVE: linalg.map outs(%[[EMPTY]] : tensor<7x10xf32>
+// CHECK-PRIMITIVE-SAME: {someattr}
+// CHECK-PRIMITIVE:        %[[INDEX:.*]] = linalg.index 1
+// CHECK-PRIMITIVE-NEXT:   %[[INT_CAST:.*]] = arith.index_cast %[[INDEX]] : index to i64
+// CHECK-PRIMITIVE-NEXT:   %[[FLOAT_CAST:.*]] = arith.sitofp %[[INT_CAST]] : i64 to f32
+// CHECK-PRIMITIVE-NEXT:   linalg.yield %[[FLOAT_CAST]]
+
+// -----
+
+// CHECK: #[[RESULT_MAP:.*]] = affine_map<(d0, d1) -> (d0, d1)>
+// CHECK-LABEL: func @iota_i32
+func.func @iota_i32() -> tensor<7x10xi32> {
+  %result = "stablehlo.iota"() {iota_dimension = 1 : i64} : () -> (tensor<7x10xi32>)
+  func.return %result : tensor<7x10xi32>
+}
+// CHECK: tensor.empty
+// CHECK: linalg.generic
+// CHECK-SAME: indexing_maps = [#[[RESULT_MAP]]]
+// CHECK-NEXT: ^bb0(%{{.*}}: i32):
+// CHECK-NEXT:   %[[INDEX:.*]] = linalg.index 1
+// CHECK-NEXT:   %[[INT_CAST:.*]] = arith.index_cast %[[INDEX]] : index to i32
+// CHECK-NEXT:   linalg.yield %[[INT_CAST]] : i32
+
+// -----
+
+// CHECK: #[[RESULT_MAP:.*]] = affine_map<(d0, d1) -> (d0, d1)>
+// CHECK-LABEL: func @iota_ui32
+func.func @iota_ui32() -> tensor<7x10xui32> {
+  %result = "stablehlo.iota"() {iota_dimension = 1 : i64} : () -> (tensor<7x10xui32>)
+  func.return %result : tensor<7x10xui32>
+}
+// CHECK: tensor.empty
+// CHECK: linalg.generic
+// CHECK-SAME: indexing_maps = [#[[RESULT_MAP]]]
+// CHECK-NEXT: ^bb0(%{{.*}}: i32):
+// CHECK-NEXT:   %[[INDEX:.*]] = linalg.index 1
+// CHECK-NEXT:   %[[INT_CAST:.*]] = arith.index_cast %[[INDEX]] : index to i32
+// CHECK-NEXT:   linalg.yield %[[INT_CAST]] : i32
+// CHECK: builtin.unrealized_conversion_cast %{{.*}} : tensor<7x10xi32> to tensor<7x10xui32>
+
+// -----
+
+// CHECK: #[[RESULT_MAP:.*]] = affine_map<(d0, d1) -> (d0, d1)>
+// CHECK-LABEL: func @iota_complexf32
+func.func @iota_complexf32() -> tensor<7x10xcomplex<f32>> {
+  %result = "stablehlo.iota"() {iota_dimension = 1 : i64} : () -> (tensor<7x10xcomplex<f32>>)
+  func.return %result : tensor<7x10xcomplex<f32>>
+}
+// CHECK-DAG:    %[[ZERO:.*]] = arith.constant 0.000000e+00 : f32
+// CHECK: tensor.empty
+// CHECK: linalg.generic
+// CHECK-SAME: indexing_maps = [#[[RESULT_MAP]]]
+// CHECK-NEXT: ^bb0(%{{.*}}: complex<f32>):
+// CHECK-NEXT:   %[[INDEX:.*]] = linalg.index 1
+// CHECK-NEXT:   %[[INT_CAST:.*]] = arith.index_cast %[[INDEX]] : index to i32
+// CHECK-NEXT:   %[[FLOAT_CAST:.*]] = arith.sitofp %[[INT_CAST]] : i32 to f32
+// CHECK-NEXT:   %[[COMPLEX_CAST:.*]] = complex.create %[[FLOAT_CAST]], %[[ZERO]] : complex<f32>
+// CHECK-NEXT:   linalg.yield %[[COMPLEX_CAST]] : complex<f32>
+
+// -----
+
+// CHECK: #[[RESULT_MAP:.*]] = affine_map<(d0, d1, d2) -> (d0, d1, d2)>
+// CHECK-LABEL: func @dynamic_iota_f32
+// CHECK-SAME: %[[SHAPE:.*]]: tensor<?xi32>
+func.func @dynamic_iota_f32(%shape: tensor<?xi32>) -> tensor<?x?x8xf32> {
+  %result = "stablehlo.dynamic_iota"(%shape) {iota_dimension = 1 : i64} : (tensor<?xi32>) -> (tensor<?x?x8xf32>)
+  func.return %result : tensor<?x?x8xf32>
+}
+// CHECK: %[[V1:.*]] = tensor.extract %[[SHAPE]][%c0]
+// CHECK: %[[I1:.*]] = arith.index_cast %[[V1]] : i32 to index
+// CHECK: %[[V2:.*]] = tensor.extract %[[SHAPE]][%c1]
+// CHECK: %[[I2:.*]] = arith.index_cast %[[V2]] : i32 to index
+// CHECK: tensor.empty(%[[I1]], %[[I2]]) : tensor<?x?x8xf32>
+// CHECK: linalg.generic
+// CHECK-SAME: indexing_maps = [#[[RESULT_MAP]]]
+// CHECK-NEXT: ^bb0(%{{.*}}: f32):
+// CHECK-NEXT:   %[[INDEX:.*]] = linalg.index 1
+// CHECK-NEXT:   %[[INT_CAST:.*]] = arith.index_cast %[[INDEX]] : index to i32
+// CHECK-NEXT:   %[[FLOAT_CAST:.*]] = arith.sitofp %[[INT_CAST]] : i32 to f32
+// CHECK-NEXT:   linalg.yield %[[FLOAT_CAST]] : f32
+
+// -----
+
+// CHECK: #[[RESULT_MAP:.*]] = affine_map<(d0, d1, d2) -> (d0, d1, d2)>
+// CHECK-LABEL: func @dyanmic_iota_ui32
+// CHECK-SAME: %[[SHAPE:.*]]: tensor<?xi32>
+func.func @dyanmic_iota_ui32(%shape: tensor<?xi32>) -> tensor<?x?x8xui32> {
+  %result = "stablehlo.dynamic_iota"(%shape) {iota_dimension = 1 : i64} : (tensor<?xi32>) -> (tensor<?x?x8xui32>)
+  func.return %result : tensor<?x?x8xui32>
+}
+// CHECK: %[[V1:.*]] = tensor.extract %[[SHAPE]][%c0]
+// CHECK: %[[I1:.*]] = arith.index_cast %[[V1]] : i32 to index
+// CHECK: %[[V2:.*]] = tensor.extract %[[SHAPE]][%c1]
+// CHECK: %[[I2:.*]] = arith.index_cast %[[V2]] : i32 to index
+// CHECK: tensor.empty(%[[I1]], %[[I2]]) : tensor<?x?x8xi32>
+// CHECK: linalg.generic
+// CHECK-SAME: indexing_maps = [#[[RESULT_MAP]]]
+// CHECK-NEXT: ^bb0(%{{.*}}: i32):
+// CHECK-NEXT:   %[[INDEX:.*]] = linalg.index 1
+// CHECK-NEXT:   %[[INT_CAST:.*]] = arith.index_cast %[[INDEX]] : index to i32
+// CHECK-NEXT:   linalg.yield %[[FLOAT_CAST]] : i32
+// CHECK: builtin.unrealized_conversion_cast %{{.*}} : tensor<?x?x8xi32> to tensor<?x?x8xui32>
+
+// -----
+
+func.func @map_mixed(%arg0: tensor<?xf32>,
+                     %arg1: tensor<4xf32>) -> tensor<?xf32> {
+  %0 = "stablehlo.map"(%arg0, %arg1) ({
+  ^bb0(%arg2: tensor<f32>, %arg3: tensor<f32>):
+    %1 = stablehlo.add %arg2, %arg3 : tensor<f32>
+    "stablehlo.return"(%1) : (tensor<f32>) -> ()
+  }) {dimensions = dense<0> : tensor<1xi64>}
+  : (tensor<?xf32>, tensor<4xf32>) -> tensor<?xf32>
+  func.return %0 : tensor<?xf32>
+}
+
+// CHECK-LABEL: @map_mixed
+// CHECK: linalg.generic
+
+// CHECK-PRIMITIVE-LABEL: @map_mixed
+// CHECK-PRIMITIVE: linalg.map
+
+// -----
+
+func.func @map_one_arg(%arg0: tensor<?xf32>) -> tensor<?xf32> {
+  %0 = "stablehlo.map"(%arg0) ({
+  ^bb0(%arg2: tensor<f32>):
+    %1 = stablehlo.add %arg2, %arg2 : tensor<f32>
+    "stablehlo.return"(%1) : (tensor<f32>) -> ()
+  }) {dimensions = dense<0> : tensor<1xi64>}
+  : (tensor<?xf32>) -> tensor<?xf32>
+  func.return %0 : tensor<?xf32>
+}
+
+// CHECK-LABEL: @map_one_arg
+// CHECK: linalg.generic
+
+// CHECK-PRIMITIVE-LABEL: @map_one_arg
+// CHECK-PRIMITIVE: linalg.map
+
+// -----
+
+// CHECK-LABEL: @reduce_add_unranked
+// CHECK-PRIMITIVE-LABEL: @reduce_add_unranked
+func.func @reduce_add_unranked(%arg0: tensor<*xi32>, %arg1: tensor<i32>) -> tensor<*xi32> {
+  %0 = "stablehlo.reduce"(%arg0, %arg1) ({
+  ^bb0(%arg3: tensor<i32>, %arg4 : tensor<i32>):
+    %1 = stablehlo.add %arg3, %arg4 : tensor<i32>
+    "stablehlo.return"(%1) : (tensor<i32>) -> ()
+  }) {dimensions = dense<1> : tensor<1xi64>, someattr} : (tensor<*xi32>, tensor<i32>) -> tensor<*xi32>
+  func.return %0 : tensor<*xi32>
+}
+// CHECK: stablehlo.reduce
+// CHECK-PRIMITIVE: stablehlo.reduce
+
+// -----
+
+func.func @reduce_dynamic_output(%arg0: tensor<5x4xi32>, %arg1: tensor<i32>) -> tensor<?xi32> {
+  %0 = "stablehlo.reduce"(%arg0, %arg1) ({
+  ^bb0(%arg3: tensor<i32>, %arg4 : tensor<i32>):
+    %1 = stablehlo.maximum %arg3, %arg4 : tensor<i32>
+    "stablehlo.return"(%1) : (tensor<i32>) -> ()
+  }) {dimensions = dense<0> : tensor<1xi64>} : (tensor<5x4xi32>, tensor<i32>) -> tensor<?xi32>
+  func.return %0 : tensor<?xi32>
+}
+
+// Regression test: just check that this lowers successfully.
+// CHECK-LABEL: @reduce_dynamic_output
+// CHECK: linalg.generic
+
+// CHECK-PRIMITIVE-LABEL: @reduce_dynamic_output
+// CHECK-PRIMITIVE: linalg.reduce
+
+// -----
+
+func.func @pad_cst(%arg0: tensor<12x4xf32>) -> tensor<18x12xf32> {
+  %0 = arith.constant dense<0.0> : tensor<f32>
+  %1 = "stablehlo.pad"(%arg0, %0) {
+    edge_padding_high = dense<[2, 3]> : tensor<2xi64>,
+    edge_padding_low = dense<[4, 5]> : tensor<2xi64>,
+    interior_padding = dense<0> : tensor<2xi64>
+  } : (tensor<12x4xf32>, tensor<f32>) -> tensor<18x12xf32>
+  func.return %1 : tensor<18x12xf32>
+}
+// CHECK-LABEL: func @pad_cst
+// CHECK-SAME:    %[[ARG0:[a-zA-Z0-9_]*]]
+//   CHECK-DAG: %[[CST:.+]] = arith.constant 0.000000e+00 : f32
+//       CHECK: tensor.pad %[[ARG0]] low[4, 5] high[2, 3]
+//       CHECK:  tensor.yield %[[CST]] : f32
+//       CHECK: } : tensor<12x4xf32> to tensor<18x12xf32>
+
+// -----
+
+func.func @pad_tensor(%arg0: tensor<12x4xf32>, %arg1: tensor<f32>) -> tensor<18x12xf32> {
+  %0 = "stablehlo.pad"(%arg0, %arg1) {
+    edge_padding_high = dense<[2, 3]> : tensor<2xi64>,
+    edge_padding_low = dense<[4, 5]> : tensor<2xi64>,
+    interior_padding = dense<0> : tensor<2xi64>
+  } : (tensor<12x4xf32>, tensor<f32>) -> tensor<18x12xf32>
+  func.return %0 : tensor<18x12xf32>
+}
+// CHECK-LABEL: func @pad_tensor
+//  CHECK-SAME:   %[[ARG0:[a-zA-Z0-9_]*]]
+//  CHECK-SAME:   %[[ARG1:[a-zA-Z0-9_]*]]
+//   CHECK-DAG:   %[[PAD:.+]] = tensor.extract %[[ARG1]][] : tensor<f32>
+//       CHECK:   tensor.pad %[[ARG0]] low[4, 5] high[2, 3]
+//       CHECK:     tensor.yield %[[PAD]] : f32
+//       CHECK:   } : tensor<12x4xf32> to tensor<18x12xf32>
+
+// -----
+
+func.func @pad_interior(%arg0: tensor<12x4xui32>, %arg1: tensor<ui32>) -> tensor<29x15xui32> {
+  %0 = arith.constant dense<0> : tensor<ui32>
+  %1 = "stablehlo.pad"(%arg0, %arg1) {
+    edge_padding_high = dense<[2, 3]> : tensor<2xi64>,
+    edge_padding_low = dense<[4, 5]> : tensor<2xi64>,
+    interior_padding = dense<[1, 1]> : tensor<2xi64>
+  } : (tensor<12x4xui32>, tensor<ui32>) -> tensor<29x15xui32>
+  func.return %1 : tensor<29x15xui32>
+}
+// CHECK-LABEL: func @pad_interior
+// CHECK-SAME:    %[[ARG0:[a-zA-Z0-9_]*]]
+//  CHECK-SAME:   %[[ARG1:[a-zA-Z0-9_]*]]
+//   CHECK-DAG: %[[CAST0:.+]] = builtin.unrealized_conversion_cast %[[ARG0]] : tensor<12x4xui32> to tensor<12x4xi32>
+//   CHECK-DAG: %[[CAST1:.+]] = builtin.unrealized_conversion_cast %[[ARG1]] : tensor<ui32> to tensor<i32>
+//   CHECK-DAG: %[[PAD:.+]] = tensor.extract %[[CAST1]][] : tensor<i32>
+//       CHECK: %[[INIT:.+]] = tensor.empty() : tensor<29x15xi32>
+//       CHECK: %[[FILL:.+]] = linalg.fill ins(%[[PAD]] : i32) outs(%[[INIT]] : tensor<29x15xi32>) -> tensor<29x15xi32>
+//       CHECK: %[[INSERT:.+]] = tensor.insert_slice %[[CAST0]] into %[[FILL]][4, 5] [12, 4] [2, 2] : tensor<12x4xi32> into tensor<29x15xi32>
+
+// -----
+
+func.func @pad_interior_negative(%arg0: tensor<12x4xui32>, %arg1: tensor<ui32>) -> tensor<25x9xui32> {
+  %0 = arith.constant dense<0> : tensor<ui32>
+  %1 = "stablehlo.pad"(%arg0, %arg1) {
+    edge_padding_high = dense<[-2, 3]> : tensor<2xi64>,
+    edge_padding_low = dense<[4, -1]> : tensor<2xi64>,
+    interior_padding = dense<[1, 1]> : tensor<2xi64>
+  } : (tensor<12x4xui32>, tensor<ui32>) -> tensor<25x9xui32>
+  func.return %1 : tensor<25x9xui32>
+}
+// CHECK-LABEL: func @pad_interior_negative
+//       CHECK: %[[PAD:.*]] = tensor.insert_slice %{{.+}} into %{{.+}}[4, 0] [12, 4] [2, 2] : tensor<12x4xi32> into tensor<29x10xi32>
+//       CHECK: %[[SLICE:.*]] = tensor.extract_slice %[[PAD]][0, 1] [25, 9] [1, 1] : tensor<29x10xi32> to tensor<25x9xi32>
+
+// -----
+
+func.func @torch_index_select(%arg0: tensor<5x1x5xi32>,
+                         %arg1: tensor<2xi32>) ->  tensor<2x1x5xi32> {
+  %0 = "stablehlo.torch_index_select"(%arg0, %arg1) {
+    dim = 0 : i64,
+    batch_dims = 0 : i64,
+    someattr
+  } : (tensor<5x1x5xi32>, tensor<2xi32>) -> tensor<2x1x5xi32>
+  func.return %0 : tensor<2x1x5xi32>
+}
+//  CHECK-DAG: #[[MAP0:.+]] = affine_map<(d0, d1, d2) -> (d0)>
+//  CHECK-DAG: #[[MAP1:.+]] = affine_map<(d0, d1, d2) -> (d1, d2)>
+//  CHECK-DAG: #[[MAP2:.+]] = affine_map<(d0, d1, d2) -> (d0, d1, d2)>
+//      CHECK: func @torch_index_select
+// CHECK-SAME:   %[[INPUT:[a-zA-Z0-9_]*]]
+// CHECK-SAME:   %[[INDEX:[a-zA-Z0-9_]*]]
+//      CHECK: %[[INIT1:.+]] = tensor.empty() :
+//      CHECK: %[[INIT2:.+]] = tensor.empty() :
+//      CHECK: linalg.generic {
+// CHECK-SAME:   indexing_maps
+// CHECK-SAME:   #[[MAP0]], #[[MAP1]], #[[MAP2]]
+// CHECK-SAME:   iterator_types = ["parallel", "parallel", "parallel"]
+// CHECK-SAME: ins(%[[INDEX]], %[[INIT1]] :
+// CHECK-SAME: outs(%[[INIT2]] :
+// CHECK-SAME: {someattr}
+//      CHECK: ^{{.+}}(%[[VAL:.+]]: i32, %{{.+}}: i32, %{{.+}}: i32):
+//      CHECK:   %[[CAST:.+]] = arith.index_cast %[[VAL]] : i32 to index
+//      CHECK:   %[[J:.+]] = linalg.index 1
+//      CHECK:   %[[K:.+]] = linalg.index 2
+//      CHECK:   %[[VAL2:.+]] = tensor.extract %[[INPUT]][%[[CAST]], %[[J]], %[[K]]] : tensor<5x1x5xi32>
+//      CHECK:   linalg.yield %[[VAL2]] : i32
+
+
+// -----
+
+func.func @dot_general(%arg0: tensor<?x?x?xf32>,
+                  %arg1: tensor<?x?x?xf32>) -> tensor<?x?x?xf32> {
+  %0 = "stablehlo.dot_general"(%arg0, %arg1) {
+    dot_dimension_numbers = #stablehlo.dot<
+      lhs_batching_dimensions = [1],
+      lhs_contracting_dimensions = [2],
+      rhs_batching_dimensions = [2],
+      rhs_contracting_dimensions = [1]
+    >,
+    precision_config = [#stablehlo<precision DEFAULT>, #stablehlo<precision DEFAULT>],
+    someattr
+  } : (tensor<?x?x?xf32>, tensor<?x?x?xf32>) -> tensor<?x?x?xf32>
+  func.return %0 : tensor<?x?x?xf32>
+}
+// The iterations are (Batch Dim, LHS Other Dim, RHS Other dim, Contracting Dim)
+// CHECK: #[[MAP0:.*]] = affine_map<(d0, d1, d2, d3) -> (d1, d0, d3)>
+// CHECK: #[[MAP1:.*]] = affine_map<(d0, d1, d2, d3) -> (d2, d3, d0)>
+// Output is the iterators excluding contracting
+// CHECK: #[[MAP2:.*]] = affine_map<(d0, d1, d2, d3) -> (d0, d1, d2)>
+// CHECK: func @dot_general(
+// CHECK-SAME: %[[ARG0:.*]]: tensor<?x?x?xf32>, %[[ARG1:.*]]: tensor<?x?x?xf32>)
+// CHECK-DAG: %[[C1:.*]] = arith.constant 1 : index
+// CHECK-DAG: %[[C0:.*]] = arith.constant 0 : index
+// CHECK: %[[D0:.*]] = tensor.dim %[[ARG0]], %[[C1]]
+// CHECK: %[[D1:.*]] = tensor.dim %[[ARG0]], %[[C0]]
+// CHECK: %[[D2:.*]] = tensor.dim %[[ARG1]], %[[C0]]
+// CHECK: %[[INIT:.*]] = tensor.empty(%[[D0]], %[[D1]], %[[D2]])
+// CHECK: %[[FILL:.*]] = linalg.fill ins(%{{.*}}{{.*}}outs(%[[INIT]]
+// CHECK: linalg.generic
+// CHECK-SAME: indexing_maps = [#[[MAP0]], #[[MAP1]], #[[MAP2]]]
+// Only contracting dims are reductions
+// CHECK-SAME: iterator_types = ["parallel", "parallel", "parallel", "reduction"]
+// CHECK-SAME: ins(%[[ARG0]], %[[ARG1]] : tensor<?x?x?xf32>, tensor<?x?x?xf32>)
+// CHECK-SAME: outs(%[[FILL]] : tensor<?x?x?xf32>)
+// CHECK-SAME: {someattr}
+// CHECK:   ^bb0(%[[ARG2:.*]]: f32, %[[ARG3:.*]]: f32, %[[ARG4:.*]]: f32):
+// CHECK:     %[[MUL:.*]] = arith.mulf %[[ARG2]], %[[ARG3]] : f32
+// CHECK:     %[[SUM:.*]] = arith.addf %[[ARG4]], %[[MUL]] : f32
+// CHECK:     linalg.yield %[[SUM]] : f32
+// CHECK: } -> tensor<?x?x?xf32>
+
+// -----
+
+func.func @dot_general_unsigned(%arg0: tensor<?x?x?xui32>,
+                  %arg1: tensor<?x?x?xui32>) -> tensor<?x?x?xui32> {
+  %0 = "stablehlo.dot_general"(%arg0, %arg1) {
+    dot_dimension_numbers = #stablehlo.dot<
+      lhs_batching_dimensions = [1],
+      lhs_contracting_dimensions = [2],
+      rhs_batching_dimensions = [2],
+      rhs_contracting_dimensions = [1]
+    >,
+    precision_config = [#stablehlo<precision DEFAULT>, #stablehlo<precision DEFAULT>],
+    someattr
+  } : (tensor<?x?x?xui32>, tensor<?x?x?xui32>) -> tensor<?x?x?xui32>
+  func.return %0 : tensor<?x?x?xui32>
+}
+
+// CHECK-LABEL: func @dot_general_unsigned(
+// CHECK: linalg.generic
+// CHECK-SAME: ins({{.*}} : tensor<?x?x?xi32>, tensor<?x?x?xi32>)
+// CHECK-SAME: outs({{.*}} : tensor<?x?x?xi32>)
+
+// -----
+
+func.func @dot_general_complex(%arg0: tensor<?x?x?xcomplex<f32>>,
+                  %arg1: tensor<?x?x?xcomplex<f32>>) -> tensor<?x?x?xcomplex<f32>> {
+  %0 = "stablehlo.dot_general"(%arg0, %arg1) {
+    dot_dimension_numbers = #stablehlo.dot<
+      lhs_batching_dimensions = [1],
+      lhs_contracting_dimensions = [2],
+      rhs_batching_dimensions = [2],
+      rhs_contracting_dimensions = [1]
+    >,
+    precision_config = [#stablehlo<precision DEFAULT>, #stablehlo<precision DEFAULT>],
+    someattr
+  } : (tensor<?x?x?xcomplex<f32>>, tensor<?x?x?xcomplex<f32>>) -> tensor<?x?x?xcomplex<f32>>
+  func.return %0 : tensor<?x?x?xcomplex<f32>>
+}
+
+// CHECK-LABEL: func @dot_general_complex(
+// CHECK: linalg.generic
+// CHECK: complex.mul
+// CHECK: complex.add
+
+// -----
+
+func.func @dot_general_multiple_batch_dimensions(%arg0: tensor<3x4x2x4xi32>,
+             %arg1: tensor<3x4x3x2xi32>) -> tensor<3x4x4x3xi32> {
+  %0 = "stablehlo.dot_general"(%arg0, %arg1) {
+    dot_dimension_numbers = #stablehlo.dot<
+      lhs_batching_dimensions = [0, 1],
+      rhs_batching_dimensions = [0, 1],
+      lhs_contracting_dimensions = [2],
+      rhs_contracting_dimensions = [3]>,
+    precision_config = [#stablehlo<precision DEFAULT>, #stablehlo<precision DEFAULT>],
+    someattr
+  } : (tensor<3x4x2x4xi32>, tensor<3x4x3x2xi32>) -> tensor<3x4x4x3xi32>
+  return %0 : tensor<3x4x4x3xi32>
+}
+
+// CHECK: #[[MAP0:.*]] = affine_map<(d0, d1, d2, d3, d4) -> (d0, d1, d4, d2)>
+// CHECK: #[[MAP1:.*]] = affine_map<(d0, d1, d2, d3, d4) -> (d0, d1, d3, d4)>
+// CHECK: #[[MAP2:.*]] = affine_map<(d0, d1, d2, d3, d4) -> (d0, d1, d2, d3)>
+// CHECK-LABEL: func @dot_general_multiple_batch_dimensions(
+// CHECK: linalg.generic
+// CHECK-SAME: indexing_maps = [#[[MAP0]], #[[MAP1]], #[[MAP2]]]
+// CHECK-SAME: iterator_types = ["parallel", "parallel", "parallel", "parallel", "reduction"]}
+// CHECK-SAME: ins(%[[ARG0]], %[[ARG1]] : tensor<3x4x2x4xi32>, tensor<3x4x3x2xi32>)
+// CHECK-SAME: outs({{.*}} : tensor<3x4x4x3xi32>)
+// CHECK-SAME: {someattr}
+
+// -----
+
+// CHECK-LABEL: set_dimension_size
+// CHECK-SAME: %[[VALUE:.*]]: tensor<2x?xf32, #stablehlo.bounds<?, 2>
+func.func @set_dimension_size(
+  %value: tensor<2x?xf32, #stablehlo.type_extensions<bounds = [?, 2]>>,
+  %dimension: tensor<i32>)
+  -> tensor<2x?xf32, #stablehlo.type_extensions<bounds = [?, 2]>> {
+  // CHECK: tensor.extract_slice %[[VALUE]][0, 0] [2, %{{.*}}] [1, 1] : tensor<2x?xf32, #stablehlo.bounds<?, 2>> to tensor<2x?xf32, #stablehlo.bounds<?, 2>>
+  %0 = "stablehlo.set_dimension_size"(%value, %dimension) { dimension = 1 }
+    : (tensor<2x?xf32, #stablehlo.type_extensions<bounds = [?, 2]>>, tensor<i32>)
+    -> tensor<2x?xf32, #stablehlo.type_extensions<bounds = [?, 2]>>
+  func.return %0 : tensor<2x?xf32, #stablehlo.type_extensions<bounds = [?, 2]>>
+}
+
+// -----
+
+// CHECK-DAG: #[[OPERAND_MAP:.*]] = affine_map<(d0, d1, d2, d3) -> (d1, d0, d3, d2)>
+// CHECK-DAG: #[[RESULT_MAP:.*]] = affine_map<(d0, d1, d2, d3) -> (d0, d1, d2, d3)>
+// CHECK-LABEL: func @transpose
+func.func @transpose(%arg0: tensor<2x3x9x5xi32>) -> tensor<3x2x5x9xi32> {
+  %0 = "stablehlo.transpose"(%arg0) {permutation = dense<[1, 0, 3, 2]> : tensor<4xi64>}
+        : (tensor<2x3x9x5xi32>) -> tensor<3x2x5x9xi32>
+  func.return %0 : tensor<3x2x5x9xi32>
+}
+// CHECK: linalg.generic {{{.*}}indexing_maps = [#[[OPERAND_MAP]], #[[RESULT_MAP]]]
+
+// CHECK-PRIMITIVE-LABEL: func @transpose
+// CHECK-PRIMITIVE: linalg.transpose
+
+// -----
+
+// CHECK-DAG: #[[OPERAND_MAP:.*]] = affine_map<(d0, d1, d2, d3) -> (d1, d0, d3, d2)>
+// CHECK-DAG: #[[RESULT_MAP:.*]] = affine_map<(d0, d1, d2, d3) -> (d0, d1, d2, d3)>
+// CHECK-LABEL: func @transpose_dynamic
+func.func @transpose_dynamic(%arg0: tensor<?x?x9x?xi32>) -> tensor<?x?x?x9xi32> {
+  %0 = "stablehlo.transpose"(%arg0) {permutation = dense<[1, 0, 3, 2]> : tensor<4xi64>, someattr}
+        : (tensor<?x?x9x?xi32>) -> tensor<?x?x?x9xi32>
+  func.return %0 : tensor<?x?x?x9xi32>
+}
+// CHECK-DAG: %[[C0:.*]] = arith.constant 0 : index
+// CHECK-DAG: %[[C1:.*]] = arith.constant 1 : index
+// CHECK-DAG: %[[C3:.*]] = arith.constant 3 : index
+// CHECK: %[[D0:.*]] = tensor.dim %arg0, %[[C0]]
+// CHECK: %[[D1:.*]] = tensor.dim %arg0, %[[C1]]
+// CHECK: %[[D3:.*]] = tensor.dim %arg0, %[[C3]]
+// CHECK: %[[INIT:.*]] = tensor.empty(%[[D1]], %[[D0]], %[[D3]]) : tensor<?x?x?x9xi32>
+// CHECK: linalg.generic {{{.*}}indexing_maps = [#[[OPERAND_MAP]], #[[RESULT_MAP]]]
+// CHECK-SAME: ins(%arg0 : tensor<?x?x9x?xi32>) outs(%[[INIT]] : tensor<?x?x?x9xi32>)
+// CHECK-SAME: {someattr}
+
+// CHECK-PRIMITIVE-LABEL: func @transpose_dynamic
+// CHECK-PRIMITIVE-DAG: %[[C0:.*]] = arith.constant 0 : index
+// CHECK-PRIMITIVE-DAG: %[[C1:.*]] = arith.constant 1 : index
+// CHECK-PRIMITIVE-DAG: %[[C3:.*]] = arith.constant 3 : index
+// CHECK-PRIMITIVE: %[[D0:.*]] = tensor.dim %arg0, %[[C0]]
+// CHECK-PRIMITIVE: %[[D1:.*]] = tensor.dim %arg0, %[[C1]]
+// CHECK-PRIMITIVE: %[[D3:.*]] = tensor.dim %arg0, %[[C3]]
+// CHECK-PRIMITIVE: %[[INIT:.*]] = tensor.empty(%[[D1]], %[[D0]], %[[D3]]) : tensor<?x?x?x9xi32>
+// CHECK-PRIMITIVE: linalg.transpose
+// CHECK-PRIMITIVE-SAME: ins(%arg0 : tensor<?x?x9x?xi32>)
+// CHECK-PRIMITIVE-SAME: outs(%[[INIT]] : tensor<?x?x?x9xi32>)
+// CHECK-PRIMITIVE-SAME: permutation = [1, 0, 3, 2]
+// CHECK-PRIMITIVE-SAME: {someattr}
+
+func.func @transpose_unsigned(%arg0: tensor<2x2xui32>) -> tensor<2x2xui32> {
+  %0 = "stablehlo.transpose"(%arg0) {
+    permutation = dense<[1, 0]> : tensor<2xi64>,
+    result_layout = dense<[0, 1]> : tensor<2xindex>
+  } : (tensor<2x2xui32>) -> tensor<2x2xui32>
+  return %0 : tensor<2x2xui32>
+}
+
+// Regression test. Just check that unsigned ints lower successfully.
+// CHECK-LABEL: func @transpose_unsigned
+// CHECK-PRIMITIVE-LABEL: func @transpose_unsigned
+
+// -----
+
+func.func @dot_matmul(%arg0: tensor<2x3xf32>,
+                 %arg1: tensor<3x?xf32>) -> tensor<2x?xf32> {
+  %0 = "stablehlo.dot"(%arg0, %arg1) {someattr}
+           : (tensor<2x3xf32>, tensor<3x?xf32>) -> tensor<2x?xf32>
+  func.return %0 : tensor<2x?xf32>
+}
+// CHECK-LABEL: func @dot_matmul(
+// CHECK-SAME: %[[ARG0:.*]]: tensor<2x3xf32>, %[[ARG1:.*]]: tensor<3x?xf32>)
+// CHECK: %[[C1:.*]] = arith.constant 1 : index
+// CHECK: %[[D1:.*]] = tensor.dim %[[ARG1]], %[[C1]]
+// CHECK: %[[INIT:.*]] = tensor.empty(%[[D1]])
+// CHECK: %[[FILL:.*]] = linalg.fill ins(%{{.*}}{{.*}}outs(%[[INIT]]
+// CHECK: linalg.matmul
+// CHECK-SAME: {someattr}
+// CHECK-SAME: ins(%[[ARG0]], %[[ARG1]] : tensor<2x3xf32>, tensor<3x?xf32>)
+// CHECK-SAME: outs(%[[FILL]] : tensor<2x?xf32>)
+
+// -----
+
+func.func @dot_matmul_complex(%arg0: tensor<2x3xcomplex<f32>>,
+                 %arg1: tensor<3x?xcomplex<f32>>) -> tensor<2x?xcomplex<f32>> {
+  %0 = "stablehlo.dot"(%arg0, %arg1) {someattr}
+           : (tensor<2x3xcomplex<f32>>, tensor<3x?xcomplex<f32>>) -> tensor<2x?xcomplex<f32>>
+  func.return %0 : tensor<2x?xcomplex<f32>>
+}
+// CHECK-LABEL: func @dot_matmul_complex(
+// CHECK-SAME: %[[ARG0:.*]]: tensor<2x3xcomplex<f32>>, %[[ARG1:.*]]: tensor<3x?xcomplex<f32>>)
+// CHECK: %[[C1:.*]] = arith.constant 1 : index
+// CHECK: %[[D1:.*]] = tensor.dim %[[ARG1]], %[[C1]]
+// CHECK: %[[INIT:.*]] = tensor.empty(%[[D1]]) : tensor<2x?x
+// CHECK: %[[FILL:.*]] = linalg.fill ins(%{{.*}}{{.*}}outs(%[[INIT]]
+// CHECK: linalg.matmul
+// CHECK-SAME: {someattr}
+// CHECK-SAME: ins(%[[ARG0]], %[[ARG1]] : tensor<2x3xcomplex<f32>>, tensor<3x?xcomplex<f32>>)
+// CHECK-SAME: outs(%[[FILL]] : tensor<2x?xcomplex<f32>>)
+
+// -----
+
+func.func @dot_matmul_i8_i8_i32(%arg0: tensor<2x3xi8>,
+                 %arg1: tensor<3x?xi8>) -> tensor<2x?xi32> {
+  %0 = "stablehlo.dot"(%arg0, %arg1) : (tensor<2x3xi8>,
+                                   tensor<3x?xi8>) -> tensor<2x?xi32>
+  func.return %0 : tensor<2x?xi32>
+}
+// CHECK-LABEL: func @dot_matmul_i8_i8_i32(
+// CHECK-SAME: %[[ARG0:.*]]: tensor<2x3xi8>, %[[ARG1:.*]]: tensor<3x?xi8>)
+// CHECK: %[[C1:.*]] = arith.constant 1 : index
+// CHECK: %[[D1:.*]] = tensor.dim %[[ARG1]], %[[C1]]
+// CHECK: %[[INIT:.*]] = tensor.empty(%[[D1]]) : tensor<2x?x
+// CHECK: %[[FILL:.*]] = linalg.fill ins(%{{.*}}{{.*}}outs(%[[INIT]]
+// CHECK: linalg.matmul
+// CHECK-SAME: ins(%[[ARG0]], %[[ARG1]] : tensor<2x3xi8>, tensor<3x?xi8>)
+// CHECK-SAME: outs(%[[FILL]] : tensor<2x?xi32>)
+
+// -----
+
+func.func @dot_matmul_i16_i16_i32(%arg0: tensor<2x3xi16>,
+                 %arg1: tensor<3x?xi16>) -> tensor<2x?xi32> {
+  %0 = "stablehlo.dot"(%arg0, %arg1) : (tensor<2x3xi16>,
+                                   tensor<3x?xi16>) -> tensor<2x?xi32>
+  func.return %0 : tensor<2x?xi32>
+}
+// CHECK-LABEL: func @dot_matmul_i16_i16_i32(
+// CHECK-SAME: %[[ARG0:.*]]: tensor<2x3xi16>, %[[ARG1:.*]]: tensor<3x?xi16>)
+// CHECK: %[[C1:.*]] = arith.constant 1 : index
+// CHECK: %[[D1:.*]] = tensor.dim %[[ARG1]], %[[C1]]
+// CHECK: %[[INIT:.*]] = tensor.empty(%[[D1]]) : tensor<2x?x
+// CHECK: %[[FILL:.*]] = linalg.fill ins(%{{.*}}{{.*}}outs(%[[INIT]]
+// CHECK: linalg.matmul
+// CHECK-SAME: ins(%[[ARG0]], %[[ARG1]] : tensor<2x3xi16>, tensor<3x?xi16>)
+// CHECK-SAME: outs(%[[FILL]] : tensor<2x?xi32>)
+
+// -----
+
+func.func @dot_matmul_i32_i32_i32(%arg0: tensor<2x3xi32>,
+                 %arg1: tensor<3x?xi32>) -> tensor<2x?xi32> {
+  %0 = "stablehlo.dot"(%arg0, %arg1) : (tensor<2x3xi32>,
+                                   tensor<3x?xi32>) -> tensor<2x?xi32>
+  func.return %0 : tensor<2x?xi32>
+}
+// CHECK-LABEL: func @dot_matmul_i32_i32_i32(
+// CHECK-SAME: %[[ARG0:.*]]: tensor<2x3xi32>, %[[ARG1:.*]]: tensor<3x?xi32>)
+// CHECK: %[[C1:.*]] = arith.constant 1 : index
+// CHECK: %[[D1:.*]] = tensor.dim %[[ARG1]], %[[C1]]
+// CHECK: %[[INIT:.*]] = tensor.empty(%[[D1]]) : tensor<2x?x
+// CHECK: %[[FILL:.*]] = linalg.fill ins(%{{.*}}{{.*}}outs(%[[INIT]]
+// CHECK: linalg.matmul
+// CHECK-SAME: ins(%[[ARG0]], %[[ARG1]] : tensor<2x3xi32>, tensor<3x?xi32>)
+// CHECK-SAME: outs(%[[FILL]] : tensor<2x?xi32>)
+
+// -----
+
+func.func @dot_matvec(%arg0: tensor<?x3xf32>,
+                 %arg1: tensor<3xf32>) -> tensor<?xf32> {
+  %0 = "stablehlo.dot"(%arg0, %arg1) : (tensor<?x3xf32>,
+                                   tensor<3xf32>) -> tensor<?xf32>
+  func.return %0 : tensor<?xf32>
+}
+// CHECK-LABEL: func @dot_matvec(
+// CHECK-SAME: %[[ARG0:.*]]: tensor<?x3xf32>, %[[ARG1:.*]]: tensor<3xf32>)
+// CHECK: %[[C0:.*]] = arith.constant 0 : index
+// CHECK: %[[D0:.*]] = tensor.dim %[[ARG0]], %[[C0]]
+// CHECK: %[[INIT:.*]] = tensor.empty(%[[D0]])
+// CHECK: %[[FILL:.*]] = linalg.fill ins(%{{.*}}{{.*}}outs(%[[INIT]]
+// CHECK: linalg.matvec
+// CHECK-SAME: ins(%[[ARG0]], %[[ARG1]] : tensor<?x3xf32>, tensor<3xf32>)
+// CHECK-SAME: outs(%[[FILL]] : tensor<?xf32>)
+
+// -----
+
+func.func @dot_vecmat(%arg0: tensor<3xf32>,
+                 %arg1: tensor<3x?xf32>) -> tensor<?xf32> {
+  %0 = "stablehlo.dot"(%arg0, %arg1) : (tensor<3xf32>,
+                                   tensor<3x?xf32>) -> tensor<?xf32>
+  func.return %0 : tensor<?xf32>
+}
+// CHECK-LABEL: func @dot_vecmat(
+// CHECK-SAME: %[[ARG0:.*]]: tensor<3xf32>, %[[ARG1:.*]]: tensor<3x?xf32>)
+// CHECK: %[[C1:.*]] = arith.constant 1 : index
+// CHECK: %[[D1:.*]] = tensor.dim %[[ARG1]], %[[C1]]
+// CHECK: %[[INIT:.*]] = tensor.empty(%[[D1]])
+// CHECK: linalg.fill ins(%{{.*}}{{.*}}outs(%[[INIT]]
+// CHECK: linalg.vecmat
+// CHECK-SAME: ins(%[[ARG0]], %[[ARG1]] : tensor<3xf32>, tensor<3x?xf32>)
+// CHECK-SAME: outs(%[[FILL]] : tensor<?xf32>)
+
+// -----
+
+func.func @dot_dot(%arg0: tensor<?xf32>,
+              %arg1: tensor<?xf32>) -> tensor<f32> {
+  %0 = "stablehlo.dot"(%arg0, %arg1) : (tensor<?xf32>, tensor<?xf32>) -> tensor<f32>
+  func.return %0 : tensor<f32>
+}
+// CHECK-LABEL: func @dot_dot(
+// CHECK-SAME: %[[ARG0:.*]]: tensor<?xf32>, %[[ARG1:.*]]: tensor<?xf32>)
+// CHECK: %[[INIT:.*]] = tensor.empty()
+// CHECK: %[[FILL:.*]] = linalg.fill ins(%{{.*}}{{.*}}outs(%[[INIT]]
+// CHECK: linalg.dot
+// CHECK-SAME: ins(%[[ARG0]], %[[ARG1]] : tensor<?xf32>, tensor<?xf32>)
+// CHECK-SAME: outs(%[[FILL]] : tensor<f32>)
+
+// -----
+
+func.func @dot_dot_unsigned(%arg0: tensor<?xui32>,
+              %arg1: tensor<?xui32>) -> tensor<ui32> {
+  %0 = "stablehlo.dot"(%arg0, %arg1) : (tensor<?xui32>, tensor<?xui32>) -> tensor<ui32>
+  func.return %0 : tensor<ui32>
+}
+// CHECK-LABEL: func @dot_dot_unsigned(
+// CHECK-SAME: %[[ARG0:.*]]: tensor<?xui32>, %[[ARG1:.*]]: tensor<?xui32>)
+// CHECK: %[[INIT:.*]] = tensor.empty()
+// CHECK: %[[FILL:.*]] = linalg.fill ins(%{{.*}}outs(%[[INIT]]
+// CHECK: linalg.dot
+// CHECK-SAME: ins(%{{.*}} : tensor<?xi32>, tensor<?xi32>)
+// CHECK-SAME: outs(%[[FILL]] : tensor<i32>)
diff --git a/compiler/src/iree/compiler/Pipelines/BUILD.bazel b/compiler/src/iree/compiler/Pipelines/BUILD.bazel
index 1cd3af13026c..d86154e187a9 100644
--- a/compiler/src/iree/compiler/Pipelines/BUILD.bazel
+++ b/compiler/src/iree/compiler/Pipelines/BUILD.bazel
@@ -18,6 +18,7 @@ iree_compiler_cc_library(
     hdrs = ["Options.h"],
     deps = [
         "//compiler/src/iree/compiler/InputConversion/MHLO",
+        "//compiler/src/iree/compiler/InputConversion/StableHLO",
         "//compiler/src/iree/compiler/InputConversion/TMTensor",
         "//compiler/src/iree/compiler/InputConversion/TOSA",
         "//compiler/src/iree/compiler/Utils",
@@ -48,6 +49,7 @@ iree_compiler_cc_library(
         "//compiler/src/iree/compiler/Dialect/VM/Transforms",
         "//compiler/src/iree/compiler/InputConversion/Common",
         "//compiler/src/iree/compiler/InputConversion/MHLO",
+        "//compiler/src/iree/compiler/InputConversion/StableHLO",
         "//compiler/src/iree/compiler/InputConversion/TMTensor",
         "//compiler/src/iree/compiler/InputConversion/TOSA",
         "//compiler/src/iree/compiler/Modules/HAL/Inline/Transforms",
diff --git a/compiler/src/iree/compiler/Pipelines/CMakeLists.txt b/compiler/src/iree/compiler/Pipelines/CMakeLists.txt
index 758056899b57..2508dba8666d 100644
--- a/compiler/src/iree/compiler/Pipelines/CMakeLists.txt
+++ b/compiler/src/iree/compiler/Pipelines/CMakeLists.txt
@@ -8,6 +8,7 @@
 set(IREE_INPUT_DEPS "")
 if(IREE_INPUT_MHLO)
   list(APPEND IREE_INPUT_DEPS iree::compiler::InputConversion::MHLO)
+  list(APPEND IREE_INPUT_DEPS iree::compiler::InputConversion::StableHLO)
 endif()
 if(IREE_INPUT_TORCH)
   list(APPEND IREE_INPUT_DEPS iree::compiler::InputConversion::TMTensor)
diff --git a/compiler/src/iree/compiler/Tools/BUILD.bazel b/compiler/src/iree/compiler/Tools/BUILD.bazel
index a3294fd66c7c..cf58f38cf22f 100644
--- a/compiler/src/iree/compiler/Tools/BUILD.bazel
+++ b/compiler/src/iree/compiler/Tools/BUILD.bazel
@@ -38,6 +38,7 @@ iree_compiler_cc_library(
     deps = [
         "//compiler/src/iree/compiler/InputConversion/Common",
         "//compiler/src/iree/compiler/InputConversion/MHLO",
+        "//compiler/src/iree/compiler/InputConversion/StableHLO",
         "//compiler/src/iree/compiler/InputConversion/TMTensor",
         "//compiler/src/iree/compiler/InputConversion/TOSA",
         "@llvm-project//mlir:ConversionPasses",
diff --git a/compiler/src/iree/compiler/Tools/CMakeLists.txt b/compiler/src/iree/compiler/Tools/CMakeLists.txt
index d5cd38786d44..ba0b417512d0 100644
--- a/compiler/src/iree/compiler/Tools/CMakeLists.txt
+++ b/compiler/src/iree/compiler/Tools/CMakeLists.txt
@@ -48,9 +48,11 @@ endif()
 set(IREE_INPUT_DEPS "")
 if(IREE_INPUT_MHLO)
   list(APPEND IREE_INPUT_DEPS iree::compiler::InputConversion::MHLO)
+  list(APPEND IREE_INPUT_DEPS iree::compiler::InputConversion::StableHLO)
   list(APPEND IREE_INPUT_DEPS tensorflow::external_mhlo_includes)
   list(APPEND IREE_INPUT_DEPS ChloOps)
   list(APPEND IREE_INPUT_DEPS MhloDialect)
+  list(APPEND IREE_INPUT_DEPS StablehloOps)
 endif()
 if(IREE_INPUT_TORCH)
   list(APPEND IREE_INPUT_DEPS iree::compiler::InputConversion::TMTensor)
diff --git a/compiler/src/iree/compiler/Tools/init_input_passes.cc b/compiler/src/iree/compiler/Tools/init_input_passes.cc
index 013beaaa7d52..fe9f41f79cfe 100644
--- a/compiler/src/iree/compiler/Tools/init_input_passes.cc
+++ b/compiler/src/iree/compiler/Tools/init_input_passes.cc
@@ -10,6 +10,7 @@
 
 #ifdef IREE_HAVE_MHLO_INPUT
 #include "iree/compiler/InputConversion/MHLO/Passes.h"
+#include "iree/compiler/InputConversion/StableHLO/Passes.h"
 #endif  // IREE_HAVE_MHLO_INPUT
 #ifdef IREE_HAVE_TORCH_INPUT
 #include "iree/compiler/InputConversion/TMTensor/Passes.h"
@@ -28,6 +29,7 @@ void registerInputPasses() {
 
 #ifdef IREE_HAVE_MHLO_INPUT
   MHLO::registerMHLOConversionPasses();
+  stablehlo::registerStableHLOConversionPasses();
 #endif  // IREE_HAVE_MHLO_INPUT
 #ifdef IREE_HAVE_TORCH_INPUT
   TMTensor::registerTMTensorConversionPasses();