Update llvm version

llvm-sha.txt

@@ -1 +1 @@
-6bbccd2516c3a843809a8303da48abce58a88855
+d58637219463924185614f18911c5f01a1c20aa9

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)

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)

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

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

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)

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

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

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)

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

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

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)

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

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

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>();
 }

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();
   }
 };