diff --git a/llvm-sha.txt b/llvm-sha.txt
index 42a902715..ed6390e43 100644
--- a/llvm-sha.txt
+++ b/llvm-sha.txt
@@ -1 +1 @@
-6bbccd2516c3a843809a8303da48abce58a88855
+d58637219463924185614f18911c5f01a1c20aa9
diff --git a/mlir/CMakeLists.txt b/mlir/CMakeLists.txt
index 4009cfffc..8e0b071d4 100644
--- a/mlir/CMakeLists.txt
+++ b/mlir/CMakeLists.txt
@@ -29,6 +29,10 @@ add_subdirectory(include/numba/Dialect/numba_util)
 add_subdirectory(include/numba/Dialect/gpu_runtime/IR)
 add_subdirectory(include/numba/Dialect/ntensor/IR)
 add_subdirectory(include/numba/Dialect/math_ext/IR)
+add_subdirectory(include/legacy/Dialect/Arith/Transforms)
+add_subdirectory(include/legacy/Dialect/Bufferization/Transforms)
+add_subdirectory(include/legacy/Dialect/Linalg/Transforms)
+add_subdirectory(include/legacy/Dialect/Tensor/Transforms)
 
 set(SOURCES_LIST
     lib/Analysis/AliasAnalysis.cpp
@@ -56,6 +60,10 @@ set(SOURCES_LIST
     lib/Dialect/numba_util/Dialect.cpp
     lib/Dialect/plier/Dialect.cpp
     lib/ExecutionEngine/ExecutionEngine.cpp
+    lib/legacy/Dialect/Arith/Transforms/Bufferize.cpp
+    lib/legacy/Dialect/Bufferization/Transforms/Bufferize.cpp
+    lib/legacy/Dialect/Linalg/Transforms/Bufferize.cpp
+    lib/legacy/Dialect/Tensor/Transforms/Bufferize.cpp
     lib/Transforms/CallLowering.cpp
     lib/Transforms/CanonicalizeReductions.cpp
     lib/Transforms/CastLowering.cpp
@@ -184,7 +192,17 @@ target_include_directories(${NUMBA_MLIR_LIB} PUBLIC
     ${PROJECT_BINARY_DIR}/numba/include
     )
 
-add_dependencies(${NUMBA_MLIR_LIB} MLIRPlierOpsIncGen MLIRNumbaUtilOpsIncGen MLIRGpuRuntimeOpsIncGen MLIRNTensorOpsIncGen MLIRMathExtOpsIncGen)
+add_dependencies(
+    ${NUMBA_MLIR_LIB}
+    MLIRPlierOpsIncGen
+    MLIRNumbaUtilOpsIncGen
+    MLIRGpuRuntimeOpsIncGen
+    MLIRNTensorOpsIncGen
+    MLIRMathExtOpsIncGen
+    MLIRArithLegacyIncGen
+    MLIRBufferizationLegacyIncGen
+    MLIRLinalgLegacyIncGen
+    MLIRTensorLegacyIncGen)
 
 add_subdirectory(tools)
 if(NUMBA_MLIR_ENABLE_TESTS)
diff --git a/mlir/include/legacy/Dialect/Arith/Transforms/CMakeLists.txt b/mlir/include/legacy/Dialect/Arith/Transforms/CMakeLists.txt
new file mode 100644
index 000000000..b9e76bb67
--- /dev/null
+++ b/mlir/include/legacy/Dialect/Arith/Transforms/CMakeLists.txt
@@ -0,0 +1,11 @@
+# SPDX-FileCopyrightText: 2023 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+include_directories(${MLIR_INCLUDE_DIRS})
+
+set(LLVM_TARGET_DEFINITIONS Passes.td)
+mlir_tablegen(Passes.h.inc -gen-pass-decls -name Transforms)
+mlir_tablegen(Transforms.capi.h.inc -gen-pass-capi-header --prefix Transforms)
+mlir_tablegen(Transforms.capi.cpp.inc -gen-pass-capi-impl --prefix Transforms)
+add_public_tablegen_target(MLIRArithLegacyIncGen)
diff --git a/mlir/include/legacy/Dialect/Arith/Transforms/Passes.h b/mlir/include/legacy/Dialect/Arith/Transforms/Passes.h
new file mode 100644
index 000000000..27e5a83d8
--- /dev/null
+++ b/mlir/include/legacy/Dialect/Arith/Transforms/Passes.h
@@ -0,0 +1,13 @@
+#pragma once
+
+#include "mlir/Pass/Pass.h"
+
+namespace mlir {
+namespace arith {
+namespace legacy {
+/// Create a pass to bufferize arith.constant ops.
+std::unique_ptr<Pass> createConstantBufferizePass(uint64_t alignment = 0);
+
+} // namespace legacy
+} // namespace arith
+} // namespace mlir
diff --git a/mlir/include/legacy/Dialect/Arith/Transforms/Passes.td b/mlir/include/legacy/Dialect/Arith/Transforms/Passes.td
new file mode 100644
index 000000000..ee839f095
--- /dev/null
+++ b/mlir/include/legacy/Dialect/Arith/Transforms/Passes.td
@@ -0,0 +1,29 @@
+//===-- Passes.td - Arith pass definition file --------*- tablegen -*-===//
+//
+// Part of the LLVM Project, 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 MLIR_DIALECT_ARITH_TRANSFORMS_LEGACY_PASSES
+#define MLIR_DIALECT_ARITH_TRANSFORMS_LEGACY_PASSES
+
+include "mlir/Pass/PassBase.td"
+
+def ArithBufferizePass : Pass<"arith-bufferize", "ModuleOp"> {
+  let summary = "Bufferize Arith dialect ops.";
+  let description = [{
+    This pass bufferizes arith dialect ops.
+
+    This pass needs to be a module pass because it inserts memref.global
+    ops into the module, which cannot be done safely from a function pass due to
+    multi-threading. Most other bufferization passes can run in parallel at
+    function granularity.
+  }];
+  let options = [
+    Option<"alignment", "alignment", "unsigned", /*default=*/"0",
+           "Create global memrefs with a specified alignment">,
+  ];
+}
+
+#endif // MLIR_DIALECT_ARITH_TRANSFORMS_LEGACY_PASSES
diff --git a/mlir/include/legacy/Dialect/Bufferization/Transforms/CMakeLists.txt b/mlir/include/legacy/Dialect/Bufferization/Transforms/CMakeLists.txt
new file mode 100644
index 000000000..e40a54bfd
--- /dev/null
+++ b/mlir/include/legacy/Dialect/Bufferization/Transforms/CMakeLists.txt
@@ -0,0 +1,11 @@
+# SPDX-FileCopyrightText: 2023 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+include_directories(${MLIR_INCLUDE_DIRS})
+
+set(LLVM_TARGET_DEFINITIONS Passes.td)
+mlir_tablegen(Passes.h.inc -gen-pass-decls -name Transforms)
+mlir_tablegen(Transforms.capi.h.inc -gen-pass-capi-header --prefix Transforms)
+mlir_tablegen(Transforms.capi.cpp.inc -gen-pass-capi-impl --prefix Transforms)
+add_public_tablegen_target(MLIRBufferizationLegacyIncGen)
diff --git a/mlir/include/legacy/Dialect/Bufferization/Transforms/Passes.h b/mlir/include/legacy/Dialect/Bufferization/Transforms/Passes.h
new file mode 100644
index 000000000..a12591bd3
--- /dev/null
+++ b/mlir/include/legacy/Dialect/Bufferization/Transforms/Passes.h
@@ -0,0 +1,13 @@
+#pragma once
+
+#include "mlir/Pass/Pass.h"
+
+namespace mlir {
+namespace bufferization {
+namespace legacy {
+/// Create a pass that bufferizes ops from the bufferization dialect.
+std::unique_ptr<Pass> createBufferizationBufferizePass();
+
+} // namespace legacy
+} // namespace bufferization
+} // namespace mlir
diff --git a/mlir/include/legacy/Dialect/Bufferization/Transforms/Passes.td b/mlir/include/legacy/Dialect/Bufferization/Transforms/Passes.td
new file mode 100644
index 000000000..a2b8f8b08
--- /dev/null
+++ b/mlir/include/legacy/Dialect/Bufferization/Transforms/Passes.td
@@ -0,0 +1,18 @@
+//===-- Passes.td - Bufferization pass definition file --------*- tablegen -*-===//
+//
+// Part of the LLVM Project, 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 MLIR_DIALECT_BUFFERIZATION_TRANSFORMS_LEGACY_PASSES
+#define MLIR_DIALECT_BUFFERIZATION_TRANSFORMS_LEGACY_PASSES
+
+include "mlir/Pass/PassBase.td"
+
+def BufferizationBufferize : Pass<"bufferization-bufferize", "func::FuncOp"> {
+  let summary = "Bufferize the `bufferization` dialect";
+  let constructor = "mlir::bufferization::createBufferizationBufferizePass()";
+}
+
+#endif // MLIR_DIALECT_BUFFERIZATION_TRANSFORMS_LEGACY_PASSES
diff --git a/mlir/include/legacy/Dialect/Linalg/Transforms/CMakeLists.txt b/mlir/include/legacy/Dialect/Linalg/Transforms/CMakeLists.txt
new file mode 100644
index 000000000..ba0ea2c35
--- /dev/null
+++ b/mlir/include/legacy/Dialect/Linalg/Transforms/CMakeLists.txt
@@ -0,0 +1,11 @@
+# SPDX-FileCopyrightText: 2023 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+include_directories(${MLIR_INCLUDE_DIRS})
+
+set(LLVM_TARGET_DEFINITIONS Passes.td)
+mlir_tablegen(Passes.h.inc -gen-pass-decls -name Transforms)
+mlir_tablegen(Transforms.capi.h.inc -gen-pass-capi-header --prefix Transforms)
+mlir_tablegen(Transforms.capi.cpp.inc -gen-pass-capi-impl --prefix Transforms)
+add_public_tablegen_target(MLIRLinalgLegacyIncGen)
diff --git a/mlir/include/legacy/Dialect/Linalg/Transforms/Passes.h b/mlir/include/legacy/Dialect/Linalg/Transforms/Passes.h
new file mode 100644
index 000000000..a462146da
--- /dev/null
+++ b/mlir/include/legacy/Dialect/Linalg/Transforms/Passes.h
@@ -0,0 +1,13 @@
+#pragma once
+
+#include "mlir/Pass/Pass.h"
+
+namespace mlir {
+namespace linalg {
+namespace legacy {
+/// Creates an instance of the `linalg` dialect bufferization pass.
+std::unique_ptr<Pass> createLinalgBufferizePass();
+
+} // namespace legacy
+} // namespace linalg
+} // namespace mlir
diff --git a/mlir/include/legacy/Dialect/Linalg/Transforms/Passes.td b/mlir/include/legacy/Dialect/Linalg/Transforms/Passes.td
new file mode 100644
index 000000000..3bf3cf39d
--- /dev/null
+++ b/mlir/include/legacy/Dialect/Linalg/Transforms/Passes.td
@@ -0,0 +1,24 @@
+//===-- Passes.td - Linalg pass definition file ------------*- tablegen -*-===//
+//
+// Part of the LLVM Project, 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 MLIR_DIALECT_LINALG_LEGACY_PASSES
+#define MLIR_DIALECT_LINALG_LEGACY_PASSES
+
+include "mlir/Pass/PassBase.td"
+
+def LinalgBufferizePass : Pass<"linalg-bufferize"> {
+  let summary = "Bufferize the linalg dialect";
+  let dependentDialects = [
+    "affine::AffineDialect",
+    "bufferization::BufferizationDialect",
+    "linalg::LinalgDialect",
+    "memref::MemRefDialect",
+  ];
+}
+
+#endif // MLIR_DIALECT_LINALG_LEGACY_PASSES
diff --git a/mlir/include/legacy/Dialect/Tensor/Transforms/CMakeLists.txt b/mlir/include/legacy/Dialect/Tensor/Transforms/CMakeLists.txt
new file mode 100644
index 000000000..74dfda685
--- /dev/null
+++ b/mlir/include/legacy/Dialect/Tensor/Transforms/CMakeLists.txt
@@ -0,0 +1,11 @@
+# SPDX-FileCopyrightText: 2023 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+include_directories(${MLIR_INCLUDE_DIRS})
+
+set(LLVM_TARGET_DEFINITIONS Passes.td)
+mlir_tablegen(Passes.h.inc -gen-pass-decls -name Transforms)
+mlir_tablegen(Transforms.capi.h.inc -gen-pass-capi-header --prefix Transforms)
+mlir_tablegen(Transforms.capi.cpp.inc -gen-pass-capi-impl --prefix Transforms)
+add_public_tablegen_target(MLIRTensorLegacyIncGen)
diff --git a/mlir/include/legacy/Dialect/Tensor/Transforms/Passes.h b/mlir/include/legacy/Dialect/Tensor/Transforms/Passes.h
new file mode 100644
index 000000000..a6092ca35
--- /dev/null
+++ b/mlir/include/legacy/Dialect/Tensor/Transforms/Passes.h
@@ -0,0 +1,13 @@
+#pragma once
+
+#include "mlir/Pass/Pass.h"
+
+namespace mlir {
+namespace tensor {
+namespace legacy {
+/// Creates an instance of the `tensor` dialect bufferization pass.
+std::unique_ptr<Pass> createTensorBufferizePass();
+
+} // namespace legacy
+} // namespace tensor
+} // namespace mlir
diff --git a/mlir/include/legacy/Dialect/Tensor/Transforms/Passes.td b/mlir/include/legacy/Dialect/Tensor/Transforms/Passes.td
new file mode 100644
index 000000000..772bcb9a5
--- /dev/null
+++ b/mlir/include/legacy/Dialect/Tensor/Transforms/Passes.td
@@ -0,0 +1,19 @@
+//===-- Passes.td - pass definition file -------------------*- tablegen -*-===//
+//
+// Part of the LLVM Project, 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 MLIR_DIALECT_TENSOR_TRANSFORMS_PASSES
+#define MLIR_DIALECT_TENSOR_TRANSFORMS_PASSES
+
+include "mlir/Pass/PassBase.td"
+
+def TensorBufferize : Pass<"tensor-bufferize", "func::FuncOp"> {
+  let summary = "Bufferize the `tensor` dialect";
+  let constructor = "mlir::tensor::createTensorBufferizePass()";
+}
+
+#endif // MLIR_DIALECT_TENSOR_TRANSFORMS_PASSES
diff --git a/mlir/lib/Conversion/GpuRuntimeToLlvm.cpp b/mlir/lib/Conversion/GpuRuntimeToLlvm.cpp
index 96f1a1b84..19ca1860b 100644
--- a/mlir/lib/Conversion/GpuRuntimeToLlvm.cpp
+++ b/mlir/lib/Conversion/GpuRuntimeToLlvm.cpp
@@ -855,6 +855,20 @@ class ConvertGpuSuggestBlockSizePattern
   }
 };
 
+class EraseGpuModuleOpPattern
+    : public mlir::OpRewritePattern<mlir::gpu::GPUModuleOp> {
+  using mlir::OpRewritePattern<mlir::gpu::GPUModuleOp>::OpRewritePattern;
+
+  mlir::LogicalResult
+  matchAndRewrite(mlir::gpu::GPUModuleOp op,
+                  mlir::PatternRewriter &rewriter) const override {
+    // GPU kernel modules are no longer necessary since we have a global
+    // constant with the CUBIN, or HSACO data.
+    rewriter.eraseOp(op);
+    return mlir::success();
+  }
+};
+
 struct GPUToLLVMPass
     : public mlir::PassWrapper<GPUToLLVMPass,
                                mlir::OperationPass<mlir::ModuleOp>> {
@@ -868,8 +882,7 @@ struct GPUToLLVMPass
 
     mlir::populateAsyncStructuralTypeConversionsAndLegality(converter, patterns,
                                                             target);
-    mlir::populateGpuToLLVMConversionPatterns(
-        converter, patterns, mlir::gpu::getDefaultGpuBinaryAnnotation());
+    mlir::populateGpuToLLVMConversionPatterns(converter, patterns);
 
     numba::populateControlFlowTypeConversionRewritesAndTarget(converter,
                                                               patterns, target);
@@ -923,6 +936,8 @@ void gpu_runtime::populateGpuToLLVMPatternsAndLegality(
       // clang-format on
       >(converter);
 
+  patterns.add<EraseGpuModuleOpPattern>(&converter.getContext());
+
   target.addIllegalDialect<mlir::gpu::GPUDialect>();
   target.addIllegalDialect<gpu_runtime::GpuRuntimeDialect>();
 }
diff --git a/mlir/lib/Transforms/MakeSignless.cpp b/mlir/lib/Transforms/MakeSignless.cpp
index 1685dcb73..eea079d26 100644
--- a/mlir/lib/Transforms/MakeSignless.cpp
+++ b/mlir/lib/Transforms/MakeSignless.cpp
@@ -262,7 +262,8 @@ struct ConvertTensorExpandShape
       return mlir::failure();
 
     rewriter.replaceOpWithNewOp<mlir::tensor::ExpandShapeOp>(
-        op, newResType, adaptor.getSrc(), adaptor.getReassociation());
+        op, newResType, adaptor.getSrc(), adaptor.getReassociation(),
+        adaptor.getOutputShape(), adaptor.getStaticOutputShape());
     return mlir::success();
   }
 };
diff --git a/mlir/lib/Transforms/ShapeIntegerRangePropagation.cpp b/mlir/lib/Transforms/ShapeIntegerRangePropagation.cpp
index 6b8e508dc..f5426b99d 100644
--- a/mlir/lib/Transforms/ShapeIntegerRangePropagation.cpp
+++ b/mlir/lib/Transforms/ShapeIntegerRangePropagation.cpp
@@ -41,10 +41,11 @@ class ShapeValue {
   ShapeValue() = default;
   ShapeValue(mlir::ShapedType shaped) : shapeRanges(std::in_place) {
     shapeRanges->reserve(shaped.getRank());
-    for (auto dim : shaped.getShape())
+    for (auto dim : shaped.getShape()) {
       shapeRanges->emplace_back(mlir::ShapedType::isDynamic(dim)
                                     ? getDefaultDimRange()
                                     : getFixedDimRange(dim));
+    }
   }
   ShapeValue(mlir::ArrayAttr attr) : shapeRanges(std::in_place) {
     shapeRanges->reserve(attr.size());
@@ -395,41 +396,88 @@ class ShapeValueAnalysis
       if (generic->getNumResults() == 0)
         return;
 
-      if (!llvm::all_of(generic.getIndexingMaps(), [](mlir::Attribute map) {
+      bool allIdentity =
+          llvm::all_of(generic.getIndexingMaps(), [](mlir::Attribute map) {
             return mlir::cast<mlir::AffineMapAttr>(map).getValue().isIdentity();
-          }))
-        return;
+          });
 
-      if (!llvm::all_of(generic.getIteratorTypes(), [](mlir::Attribute map) {
+      bool allParallel =
+          llvm::all_of(generic.getIteratorTypes(), [](mlir::Attribute map) {
             return mlir::cast<mlir::linalg::IteratorTypeAttr>(map).getValue() ==
                    mlir::utils::IteratorType::parallel;
-          }))
+          });
 
-        return;
+      auto inputsNum = generic.getInputs().size();
+      auto outs = generic.getOutputs();
 
-      ShapeValue newVal;
-      for (auto arg : op->getOperands())
-        newVal = ShapeValue::intersect(
-            newVal, {mlir::cast<mlir::ShapedType>(arg.getType())});
+      mlir::SmallVector<ShapeValue> newVals(outs.size());
 
-      for (auto result : op->getResults())
-        newVal = ShapeValue::intersect(
-            newVal, {mlir::cast<mlir::ShapedType>(result.getType())});
+      for (size_t i = 0; i < outs.size(); ++i) {
+        auto resOperandShape = operands[inputsNum + i]->getValue();
 
-      for (auto input : operands) {
-        auto inpuLatticeVal = input->getValue();
-        if (inpuLatticeVal.isUninitialized())
-          return;
+        if (!resOperandShape.isUninitialized()) {
+          newVals[i] = resOperandShape;
+        } else if (auto shaped =
+                       mlir::dyn_cast<mlir::ShapedType>(outs[i].getType())) {
+          newVals[i] = ShapeValue(shaped);
+        }
+      }
+
+      if (allIdentity && allParallel) {
+        // if all indexing maps are identity all outputs must be of the same
+        // shape
+        ShapeValue newVal = [&]() {
+          for (auto &&v : newVals)
+            if (!v.isUninitialized())
+              return v;
+
+          return ShapeValue();
+        }();
+
+        if (!newVal.isUninitialized()) {
+          for (auto &&v : newVals)
+            if (!v.isUninitialized())
+              newVal = ShapeValue::intersect(newVal, v);
+        }
 
-        newVal = ShapeValue::intersect(newVal, inpuLatticeVal);
+        for (auto arg : op->getOperands())
+          newVal = ShapeValue::intersect(
+              newVal, {mlir::cast<mlir::ShapedType>(arg.getType())});
+
+        for (auto result : op->getResults())
+          newVal = ShapeValue::intersect(
+              newVal, {mlir::cast<mlir::ShapedType>(result.getType())});
+
+        for (auto input : operands) {
+          auto inputLatticeVal = input->getValue();
+          if (inputLatticeVal.isUninitialized())
+            return;
+
+          newVal = ShapeValue::intersect(newVal, inputLatticeVal);
+        }
+
+        // Since all indexing maps are identity we can propagate single result
+        // to all outputs
+        if (!newVal.isUninitialized()) {
+          for (auto &v : newVals)
+            v = ShapeValue::intersect(v, newVal);
+        }
       }
 
-      LLVM_DEBUG(llvm::dbgs() << "ShapeValueAnalysis: Shaped linalg generic: "
-                              << newVal << "\n");
+      auto debug_msg = [&]() {
+        llvm::dbgs() << "ShapeValueAnalysis: Shaped linalg generic: ";
+        for (auto &&v : newVals)
+          llvm::dbgs() << v << " ";
+
+        llvm::dbgs() << "\n";
+      };
 
-      for (auto resultLattice : results) {
-        auto changed = resultLattice->join(newVal);
-        propagateIfChanged(resultLattice, changed);
+      LLVM_DEBUG(debug_msg());
+
+      assert(results.size() == newVals.size());
+      for (size_t i = 0; i < results.size(); ++i) {
+        auto changed = results[i]->join(newVals[i]);
+        propagateIfChanged(results[i], changed);
       }
       return;
     }
@@ -581,7 +629,6 @@ class IntegerRangeAnalysisEx : public mlir::dataflow::IntegerRangeAnalysis {
           return std::nullopt;
 
         auto shape = shapeVal.getShape();
-
         auto indexVal = *index;
         if (indexVal < 0 || indexVal >= static_cast<int64_t>(shape.size()))
           return std::nullopt;
@@ -592,21 +639,53 @@ class IntegerRangeAnalysisEx : public mlir::dataflow::IntegerRangeAnalysis {
       if (newRange) {
         LLVM_DEBUG(llvm::dbgs() << "IntegerRangeAnalysisEx: New dim val: "
                                 << newRange << "\n");
-        auto changed =
-            lattice->join(mlir::dataflow::IntegerValueRange{newRange});
+
+        auto changed = lattice->join(mlir::IntegerValueRange{newRange});
         propagateIfChanged(lattice, changed);
       }
       return;
     }
 
+    if (auto select = mlir::dyn_cast<mlir::arith::SelectOp>(op)) {
+      assert(op->getNumResults() == 1);
+      assert(results.size() == 1);
+
+      auto *lattice = results.front();
+      auto newRange = [&]() {
+        auto &condLattice = operands[0]->getValue();
+        std::optional<mlir::APInt> mbCondVal =
+            condLattice.isUninitialized()
+                ? std::nullopt
+                : condLattice.getValue().getConstantValue();
+
+        const auto &trueCase = operands[1]->getValue();
+        const auto &falseCase = operands[2]->getValue();
+
+        if (mbCondVal) {
+          if (mbCondVal->isZero())
+            return falseCase;
+          else
+            return trueCase;
+        }
+
+        if (trueCase.isUninitialized() || falseCase.isUninitialized())
+          return mlir::IntegerValueRange{};
+
+        return mlir::IntegerValueRange::join(trueCase, falseCase);
+      }();
+
+      auto changed = lattice->join(mlir::IntegerValueRange{newRange});
+      propagateIfChanged(lattice, changed);
+      return;
+    }
+
     // TODO: upstream
     if (!mlir::isa<mlir::InferIntRangeInterface>(op)) {
       for (auto &&[lattice, value] : llvm::zip(results, op->getResults())) {
         if (value.getType().isIntOrIndex()) {
           propagateIfChanged(
               lattice,
-              lattice->join(
-                  mlir::dataflow::IntegerValueRange::getMaxRange(value)));
+              lattice->join(mlir::IntegerValueRange::getMaxRange(value)));
         }
       }
       return setAllToEntryStates(results);
@@ -614,6 +693,29 @@ class IntegerRangeAnalysisEx : public mlir::dataflow::IntegerRangeAnalysis {
 
     mlir::dataflow::IntegerRangeAnalysis::visitOperation(op, operands, results);
   }
+
+  void visitNonControlFlowArguments(
+      mlir::Operation *op, const mlir::RegionSuccessor &successor,
+      mlir::ArrayRef<mlir::dataflow::IntegerValueRangeLattice *> argLattices,
+      unsigned firstIndex) override {
+
+    if (auto loop = mlir::dyn_cast<mlir::LoopLikeOpInterface>(op)) {
+      if (op->getNumResults() == 0) {
+        std::optional<mlir::Value> iv = loop.getSingleInductionVar();
+        if (iv) {
+          mlir::dataflow::IntegerValueRangeLattice *ivEntry =
+              getLatticeElement(*iv);
+          propagateIfChanged(
+              ivEntry,
+              ivEntry->join(mlir::IntegerValueRange::getMaxRange(*iv)));
+          return;
+        }
+      }
+    }
+
+    mlir::dataflow::IntegerRangeAnalysis::visitNonControlFlowArguments(
+        op, successor, argLattices, firstIndex);
+  }
 };
 
 static void printShapeAnalysisState(mlir::DataFlowSolver &solver,
diff --git a/mlir/lib/legacy/Dialect/Arith/Transforms/Bufferize.cpp b/mlir/lib/legacy/Dialect/Arith/Transforms/Bufferize.cpp
new file mode 100644
index 000000000..71dd71a34
--- /dev/null
+++ b/mlir/lib/legacy/Dialect/Arith/Transforms/Bufferize.cpp
@@ -0,0 +1,68 @@
+//===- Bufferize.cpp - Bufferization for Arith ops ---------*- C++ -*-===//
+//
+// Part of the LLVM Project, 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 "mlir/Dialect/Bufferization/Transforms/Bufferize.h"
+#include "mlir/Dialect/Arith/IR/Arith.h"
+#include "mlir/Dialect/Arith/Transforms/BufferizableOpInterfaceImpl.h"
+#include "mlir/Dialect/Bufferization/IR/BufferizableOpInterface.h"
+#include "mlir/Dialect/Bufferization/IR/Bufferization.h"
+#include "mlir/Dialect/MemRef/IR/MemRef.h"
+
+#include "legacy/Dialect/Arith/Transforms/Passes.h"
+
+namespace mlir {
+namespace arith {
+#define GEN_PASS_DECL
+#define GEN_PASS_DEF_ARITHBUFFERIZEPASS
+#include "legacy/Dialect/Arith/Transforms/Passes.h.inc"
+} // namespace arith
+} // namespace mlir
+
+using namespace mlir;
+using namespace bufferization;
+
+namespace {
+/// Pass to bufferize Arith ops.
+struct ArithBufferizePass
+    : public arith::impl::ArithBufferizePassBase<ArithBufferizePass> {
+  using ArithBufferizePassBase::ArithBufferizePassBase;
+
+  ArithBufferizePass(uint64_t alignment = 0, bool constantOpOnly = false)
+      : constantOpOnly(constantOpOnly) {
+    this->alignment = alignment;
+  }
+
+  void runOnOperation() override {
+    BufferizationOptions options = getPartialBufferizationOptions();
+    if (constantOpOnly) {
+      options.opFilter.allowOperation<arith::ConstantOp>();
+    } else {
+      options.opFilter.allowDialect<arith::ArithDialect>();
+    }
+    options.bufferAlignment = alignment;
+
+    if (failed(bufferizeOp(getOperation(), options)))
+      signalPassFailure();
+  }
+
+  void getDependentDialects(DialectRegistry &registry) const override {
+    registry.insert<bufferization::BufferizationDialect, memref::MemRefDialect,
+                    arith::ArithDialect>();
+    arith::registerBufferizableOpInterfaceExternalModels(registry);
+  }
+
+private:
+  bool constantOpOnly;
+};
+} // namespace
+
+std::unique_ptr<Pass>
+mlir::arith::legacy::createConstantBufferizePass(uint64_t alignment) {
+  return std::make_unique<ArithBufferizePass>(alignment,
+                                              /*constantOpOnly=*/true);
+}
diff --git a/mlir/lib/legacy/Dialect/Bufferization/Transforms/Bufferize.cpp b/mlir/lib/legacy/Dialect/Bufferization/Transforms/Bufferize.cpp
new file mode 100644
index 000000000..09d9c9eb3
--- /dev/null
+++ b/mlir/lib/legacy/Dialect/Bufferization/Transforms/Bufferize.cpp
@@ -0,0 +1,59 @@
+//===- Bufferize.cpp - Bufferization for `tensor` dialect ops -------------===//
+//
+// Part of the LLVM Project, 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
+//
+//===----------------------------------------------------------------------===//
+//
+// This file implements bufferization of `tensor` dialect ops
+//
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Dialect/Bufferization/Transforms/Bufferize.h"
+#include "mlir/Dialect/Arith/IR/Arith.h"
+#include "mlir/Dialect/Bufferization/IR/BufferizableOpInterface.h"
+#include "mlir/Dialect/Bufferization/IR/Bufferization.h"
+#include "mlir/Dialect/Func/IR/FuncOps.h"
+#include "mlir/Dialect/MemRef/IR/MemRef.h"
+#include "mlir/Dialect/SCF/IR/SCF.h"
+#include "mlir/Dialect/Tensor/IR/Tensor.h"
+#include "mlir/Dialect/Tensor/Transforms/BufferizableOpInterfaceImpl.h"
+#include "mlir/IR/ImplicitLocOpBuilder.h"
+#include "mlir/Transforms/DialectConversion.h"
+
+#include "legacy/Dialect/Bufferization/Transforms/Passes.h"
+
+namespace mlir {
+namespace bufferization {
+#define GEN_PASS_DECL
+#define GEN_PASS_DEF_BUFFERIZATIONBUFFERIZE
+#include "legacy/Dialect/Bufferization/Transforms/Passes.h.inc"
+} // namespace bufferization
+} // namespace mlir
+
+using namespace mlir::bufferization;
+
+namespace {
+struct BufferizationBufferizePass
+    : public mlir::bufferization::impl::BufferizationBufferizeBase<
+          BufferizationBufferizePass> {
+  void runOnOperation() override {
+    BufferizationOptions options = getPartialBufferizationOptions();
+    options.opFilter.allowDialect<BufferizationDialect>();
+
+    if (failed(bufferizeOp(getOperation(), options)))
+      signalPassFailure();
+  }
+
+  void getDependentDialects(mlir::DialectRegistry &registry) const override {
+    registry.insert<mlir::bufferization::BufferizationDialect,
+                    mlir::memref::MemRefDialect>();
+  }
+};
+} // namespace
+
+std::unique_ptr<mlir::Pass>
+mlir::bufferization::legacy::createBufferizationBufferizePass() {
+  return std::make_unique<BufferizationBufferizePass>();
+}
diff --git a/mlir/lib/legacy/Dialect/Linalg/Transforms/Bufferize.cpp b/mlir/lib/legacy/Dialect/Linalg/Transforms/Bufferize.cpp
new file mode 100644
index 000000000..477867f4e
--- /dev/null
+++ b/mlir/lib/legacy/Dialect/Linalg/Transforms/Bufferize.cpp
@@ -0,0 +1,59 @@
+//===- Bufferize.cpp - Bufferization of linalg ops ------------------------===//
+//
+// Part of the LLVM Project, 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 "mlir/Dialect/Linalg/Passes.h"
+
+#include "mlir/Dialect/Affine/IR/AffineOps.h"
+#include "mlir/Dialect/Bufferization/IR/Bufferization.h"
+#include "mlir/Dialect/Bufferization/Transforms/Bufferize.h"
+#include "mlir/Dialect/Func/IR/FuncOps.h"
+#include "mlir/Dialect/Linalg/IR/Linalg.h"
+#include "mlir/Dialect/Linalg/Transforms/BufferizableOpInterfaceImpl.h"
+#include "mlir/Dialect/Linalg/Transforms/Transforms.h"
+#include "mlir/Dialect/Tensor/IR/Tensor.h"
+#include "mlir/IR/BuiltinDialect.h"
+#include "mlir/IR/Operation.h"
+#include "mlir/Pass/Pass.h"
+
+#include "legacy/Dialect/Linalg/Transforms/Passes.h"
+
+namespace mlir {
+#define GEN_PASS_DECL
+#define GEN_PASS_DEF_LINALGBUFFERIZEPASS
+#include "legacy/Dialect/Linalg/Transforms/Passes.h.inc"
+} // namespace mlir
+
+using namespace mlir;
+using namespace bufferization;
+
+namespace {
+/// Converts Linalg operations that work on tensor-type operands or results to
+/// work on buffers.
+struct LinalgBufferizePass
+    : public impl::LinalgBufferizePassBase<LinalgBufferizePass> {
+  using impl::LinalgBufferizePassBase<
+      LinalgBufferizePass>::LinalgBufferizePassBase;
+  void runOnOperation() override {
+    BufferizationOptions options = getPartialBufferizationOptions();
+    options.opFilter.allowDialect<linalg::LinalgDialect>();
+
+    if (failed(bufferizeOp(getOperation(), options)))
+      signalPassFailure();
+  }
+
+  void getDependentDialects(DialectRegistry &registry) const override {
+    registry.insert<bufferization::BufferizationDialect, memref::MemRefDialect,
+                    tensor::TensorDialect, linalg::LinalgDialect>();
+    linalg::registerBufferizableOpInterfaceExternalModels(registry);
+  }
+};
+} // namespace
+
+std::unique_ptr<Pass> mlir::linalg::legacy::createLinalgBufferizePass() {
+  return mlir::createLinalgBufferizePass();
+}
diff --git a/mlir/lib/legacy/Dialect/Tensor/Transforms/Bufferize.cpp b/mlir/lib/legacy/Dialect/Tensor/Transforms/Bufferize.cpp
new file mode 100644
index 000000000..2e0c9cf76
--- /dev/null
+++ b/mlir/lib/legacy/Dialect/Tensor/Transforms/Bufferize.cpp
@@ -0,0 +1,60 @@
+//===- Bufferize.cpp - Bufferization for `tensor` dialect ops -------------===//
+//
+// Part of the LLVM Project, 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
+//
+//===----------------------------------------------------------------------===//
+//
+// This file implements bufferization of `tensor` dialect ops
+//
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Dialect/Bufferization/Transforms/Bufferize.h"
+#include "mlir/Dialect/Arith/IR/Arith.h"
+#include "mlir/Dialect/Bufferization/IR/BufferizableOpInterface.h"
+#include "mlir/Dialect/Bufferization/IR/Bufferization.h"
+#include "mlir/Dialect/Func/IR/FuncOps.h"
+#include "mlir/Dialect/MemRef/IR/MemRef.h"
+#include "mlir/Dialect/SCF/IR/SCF.h"
+#include "mlir/Dialect/Tensor/IR/Tensor.h"
+#include "mlir/Dialect/Tensor/Transforms/BufferizableOpInterfaceImpl.h"
+#include "mlir/IR/ImplicitLocOpBuilder.h"
+#include "mlir/Transforms/DialectConversion.h"
+
+#include "legacy/Dialect/Tensor/Transforms/Passes.h"
+
+namespace mlir {
+namespace tensor {
+#define GEN_PASS_DECL
+#define GEN_PASS_DEF_TENSORBUFFERIZE
+#include "legacy/Dialect/Tensor/Transforms/Passes.h.inc"
+} // namespace tensor
+} // namespace mlir
+
+using namespace mlir;
+using namespace bufferization;
+
+namespace {
+struct TensorBufferizePass
+    : public tensor::impl::TensorBufferizeBase<TensorBufferizePass> {
+  void runOnOperation() override {
+    BufferizationOptions options = getPartialBufferizationOptions();
+    options.opFilter.allowDialect<tensor::TensorDialect>();
+
+    if (failed(bufferizeOp(getOperation(), options)))
+      signalPassFailure();
+  }
+
+  void getDependentDialects(DialectRegistry &registry) const override {
+    registry
+        .insert<bufferization::BufferizationDialect, memref::MemRefDialect,
+                tensor::TensorDialect, scf::SCFDialect, arith::ArithDialect>();
+    tensor::registerBufferizableOpInterfaceExternalModels(registry);
+  }
+};
+} // namespace
+
+std::unique_ptr<mlir::Pass> mlir::tensor::legacy::createTensorBufferizePass() {
+  return std::make_unique<TensorBufferizePass>();
+}
diff --git a/mlir/tools/level_zero_runner/LevelZeroRunner.cpp b/mlir/tools/level_zero_runner/LevelZeroRunner.cpp
index e94149086..31b8664c7 100644
--- a/mlir/tools/level_zero_runner/LevelZeroRunner.cpp
+++ b/mlir/tools/level_zero_runner/LevelZeroRunner.cpp
@@ -29,6 +29,11 @@
 #include "numba/Conversion/GpuToGpuRuntime.hpp"
 #include "numba/Conversion/UtilToLlvm.hpp"
 
+#include "legacy/Dialect/Arith/Transforms/Passes.h"
+#include "legacy/Dialect/Bufferization/Transforms/Passes.h"
+#include "legacy/Dialect/Linalg/Transforms/Passes.h"
+#include "legacy/Dialect/Tensor/Transforms/Passes.h"
+
 using namespace mlir;
 
 static LogicalResult runMLIRPasses(mlir::Operation *op,
@@ -38,15 +43,16 @@ static LogicalResult runMLIRPasses(mlir::Operation *op,
   if (failed(applyPassManagerCLOptions(passManager)))
     return mlir::failure();
 
-  passManager.addPass(arith::createConstantBufferizePass());
+  passManager.addPass(mlir::arith::legacy::createConstantBufferizePass());
   passManager.addNestedPass<mlir::func::FuncOp>(createSCFBufferizePass());
   passManager.addNestedPass<mlir::func::FuncOp>(
       bufferization::createEmptyTensorToAllocTensorPass());
-  passManager.addNestedPass<mlir::func::FuncOp>(createLinalgBufferizePass());
   passManager.addNestedPass<mlir::func::FuncOp>(
-      bufferization::createBufferizationBufferizePass());
+      mlir::linalg::legacy::createLinalgBufferizePass());
+  passManager.addNestedPass<mlir::func::FuncOp>(
+      bufferization::legacy::createBufferizationBufferizePass());
   passManager.addNestedPass<mlir::func::FuncOp>(
-      tensor::createTensorBufferizePass());
+      mlir::tensor::legacy::createTensorBufferizePass());
   passManager.addPass(func::createFuncBufferizePass());
   passManager.addNestedPass<mlir::func::FuncOp>(
       bufferization::createFinalizingBufferizePass());
@@ -59,7 +65,7 @@ static LogicalResult runMLIRPasses(mlir::Operation *op,
   passManager.addNestedPass<mlir::func::FuncOp>(createParallelLoopToGpuPass());
 
   //  passManager.addNestedPass<mlir::func::FuncOp>(
-  //      gpu_runtime::createInsertGPUAllocsPass());Z
+  //      gpu_runtime::createInsertGPUAllocsPass());
   passManager.addPass(mlir::createCanonicalizerPass());
   passManager.addNestedPass<mlir::func::FuncOp>(
       mlir::createGpuDecomposeMemrefsPass());
diff --git a/numba_mlir/numba_mlir/mlir_compiler/lib/pipelines/LowerToGpu.cpp b/numba_mlir/numba_mlir/mlir_compiler/lib/pipelines/LowerToGpu.cpp
index 557de23d8..d73ca1334 100644
--- a/numba_mlir/numba_mlir/mlir_compiler/lib/pipelines/LowerToGpu.cpp
+++ b/numba_mlir/numba_mlir/mlir_compiler/lib/pipelines/LowerToGpu.cpp
@@ -2042,8 +2042,7 @@ struct GPUToLLVMPass
 
     mlir::populateAsyncStructuralTypeConversionsAndLegality(converter, patterns,
                                                             target);
-    mlir::populateGpuToLLVMConversionPatterns(
-        converter, patterns, mlir::gpu::getDefaultGpuBinaryAnnotation());
+    mlir::populateGpuToLLVMConversionPatterns(converter, patterns);
 
     numba::populateControlFlowTypeConversionRewritesAndTarget(converter,
                                                               patterns, target);
diff --git a/numba_mlir/numba_mlir/mlir_compiler/lib/pipelines/LowerToLlvm.cpp b/numba_mlir/numba_mlir/mlir_compiler/lib/pipelines/LowerToLlvm.cpp
index 0f8948bb1..5d7c79ad3 100644
--- a/numba_mlir/numba_mlir/mlir_compiler/lib/pipelines/LowerToLlvm.cpp
+++ b/numba_mlir/numba_mlir/mlir_compiler/lib/pipelines/LowerToLlvm.cpp
@@ -353,8 +353,8 @@ static mlir::Value divStrides(mlir::Location loc, mlir::OpBuilder &builder,
                               mlir::Value strides, mlir::Value m) {
   auto arrayType = strides.getType().cast<mlir::LLVM::LLVMArrayType>();
   mlir::Value array = builder.create<mlir::LLVM::UndefOp>(loc, arrayType);
-  auto count = arrayType.getNumElements();
-  for (auto i : llvm::seq(0u, count)) {
+  size_t count = arrayType.getNumElements();
+  for (auto i : llvm::seq(size_t(0), count)) {
     mlir::Value prev = builder.create<mlir::LLVM::ExtractValueOp>(
         loc, arrayType.getElementType(), strides, i);
     mlir::Value val = builder.create<mlir::LLVM::SDivOp>(loc, prev, m);
@@ -367,8 +367,8 @@ static mlir::Value mulStrides(mlir::Location loc, mlir::OpBuilder &builder,
                               mlir::Value strides, mlir::Value m) {
   auto arrayType = strides.getType().cast<mlir::LLVM::LLVMArrayType>();
   mlir::Value array = builder.create<mlir::LLVM::UndefOp>(loc, arrayType);
-  auto count = arrayType.getNumElements();
-  for (auto i : llvm::seq(0u, count)) {
+  size_t count = arrayType.getNumElements();
+  for (auto i : llvm::seq(size_t(0), count)) {
     mlir::Value prev = builder.create<mlir::LLVM::ExtractValueOp>(
         loc, arrayType.getElementType(), strides, i);
     mlir::Value val = builder.create<mlir::LLVM::MulOp>(loc, prev, m);
diff --git a/numba_mlir/numba_mlir/mlir_compiler/lib/pipelines/PlierToLinalg.cpp b/numba_mlir/numba_mlir/mlir_compiler/lib/pipelines/PlierToLinalg.cpp
index dbdc88a1d..6dc4ab50d 100644
--- a/numba_mlir/numba_mlir/mlir_compiler/lib/pipelines/PlierToLinalg.cpp
+++ b/numba_mlir/numba_mlir/mlir_compiler/lib/pipelines/PlierToLinalg.cpp
@@ -70,6 +70,10 @@
 #include "numba/Transforms/TypeConversion.hpp"
 #include "numba/Transforms/UpliftMath.hpp"
 
+#include "legacy/Dialect/Arith/Transforms/Passes.h"
+#include "legacy/Dialect/Linalg/Transforms/Passes.h"
+#include "legacy/Dialect/Tensor/Transforms/Passes.h"
+
 #include "BasePipeline.hpp"
 #include "NumpyResolver.hpp"
 #include "PyLinalgResolver.hpp"
@@ -2058,7 +2062,7 @@ rewritePrimitiveFunc(mlir::PatternRewriter &rewriter, mlir::Location loc,
       return mlir::isa<mlir::NoneType>(val.getType());
     };
 
-    if (opName.startswith("array.") && args.size() >= 1 &&
+    if (opName.starts_with("array.") && args.size() >= 1 &&
         llvm::all_of(args.drop_front(), isNone))
       return resolver.rewriteAttr(opName, loc, rewriter, args.front());
 
@@ -4488,7 +4492,7 @@ static void populatePlierToLinalgOptPipeline(mlir::OpPassManager &pm) {
 
   pm.addPass(mlir::createReconcileUnrealizedCastsPass());
 
-  pm.addPass(mlir::arith::createConstantBufferizePass());
+  pm.addPass(mlir::arith::legacy::createConstantBufferizePass());
   pm.addNestedPass<mlir::func::FuncOp>(
       mlir::bufferization::createEmptyTensorToAllocTensorPass());
   pm.addPass(mlir::createCanonicalizerPass());
@@ -4497,9 +4501,10 @@ static void populatePlierToLinalgOptPipeline(mlir::OpPassManager &pm) {
       std::make_unique<MixedGenericsAliasAnalysis>());
   pm.addNestedPass<mlir::func::FuncOp>(std::make_unique<AdditionalBufferize>());
   pm.addNestedPass<mlir::func::FuncOp>(mlir::createSCFBufferizePass());
-  pm.addNestedPass<mlir::func::FuncOp>(mlir::createLinalgBufferizePass());
   pm.addNestedPass<mlir::func::FuncOp>(
-      mlir::tensor::createTensorBufferizePass());
+      mlir::linalg::legacy::createLinalgBufferizePass());
+  pm.addNestedPass<mlir::func::FuncOp>(
+      mlir::tensor::legacy::createTensorBufferizePass());
   pm.addPass(mlir::func::createFuncBufferizePass());
   pm.addNestedPass<mlir::func::FuncOp>(mlir::createCanonicalizerPass());
   pm.addNestedPass<mlir::func::FuncOp>(