Integrate LLVM-Project @ c8b5d30f707757a4fe4d9d0bb01f762665f6942f

Upgrades LLVM-Project base to c8b5d30f707757a4fe4d9d0bb01f762665f6942f (2024-08-12). Upgrades StableHlo base to 24d1807a9a3e0df81103a0be9be7ad28ee34c85a (v1.5.2, 2024-08-12) Also updates build instructions. Signed-off-by: Christopher Bate <[email protected]>
NVIDIA · Aug 15, 2024 · 0bb2b0b · 0bb2b0b
1 parent e8462e7
commit 0bb2b0b
Show file tree

Hide file tree

Showing 12 changed files with 192 additions and 83 deletions.
diff --git a/mlir-tensorrt/CMakeLists.txt b/mlir-tensorrt/CMakeLists.txt
@@ -31,6 +31,7 @@ include(build_tools/cmake/CompilationOptions.cmake)
 include(build_tools/cmake/Targets.cmake)
 include(build_tools/cmake/Dependencies.cmake)
 include(build_tools/cmake/ManagedLLVM.cmake)
+include(build_tools/cmake/LLVMCommit.cmake)
 include(Version.cmake)
 
 # -------------------------------------------------
@@ -59,7 +60,7 @@ set(MLIR_TRT_DOWNLOAD_TENSORRT_VERSION "10.0" CACHE STRING
 set(MLIR_TRT_PACKAGE_CACHE_DIR "" CACHE STRING "Directory where to cache downloaded C++ packages")
 set(MLIR_TRT_USE_LINKER "" CACHE STRING "Specify a linker to use (e.g. LLD); this is just an alias for LLVM_USE_LINKER")
 set(MLIR_TRT_LLVM_COMMIT
-  "99bb9a719cec9513e72ad275c1c0302b76b6c408"
+  "${LLVM_COMMIT}"
   CACHE STRING
   "LLVM source archive commit for managed llvm-project build")
 
@@ -121,7 +122,21 @@ if(PROJECT_IS_TOP_LEVEL)
     set(LLVM_RUNTIME_OUTPUT_INTDIR ${CMAKE_BINARY_DIR}/bin)
     set(LLVM_LIBRARY_OUTPUT_INTDIR ${CMAKE_BINARY_DIR}/lib)
   elseif(MLIR_TRT_LLVM_COMMIT)
-    mtrt_llvm_project()
+    mtrt_llvm_project(
+      NAME llvm_project
+      VERSION 0.0.20240812
+      URL "https://github.com/llvm/llvm-project/archive/${MLIR_TRT_LLVM_COMMIT}.zip"
+      EXCLUDE_FROM_ALL TRUE
+      SOURCE_SUBDIR "llvm"
+      PATCHES "${CMAKE_SOURCE_DIR}/build_tools/llvm-project.patch"
+
+      OPTIONS
+        "LLVM_ENABLE_PROJECTS mlir"
+        "MLIR_ENABLE_BINDINGS_PYTHON ${MLIR_TRT_ENABLE_PYTHON}"
+        "LLVM_TARGETS_TO_BUILD host"
+        "LLVM_ENABLE_ASSERTIONS ${MLIR_TRT_ENABLE_ASSERTIONS}"
+        "LLVM_USE_LINKER ${MLIR_TRT_USE_LINKER}"
+    )
   endif()
 
   # For whatever reason, StableHLO uses this CMake variable to decide whether to build
@@ -162,12 +177,13 @@ include(HandleLLVMOptions)
 # Dependencies
 # -------------------------------------------------
 
-if(MLIR_TRT_ENABLE_HLO)
+# Download Stablehlo if it isn't provided by a parent project.
+if(MLIR_TRT_ENABLE_HLO AND NOT TARGET StablehloOps)
   mtrt_add_stablehlo(
     # We don't tag versions vX.Y.Z in stablehlo, so use the upgrade
     # year-month-day as a tag for CPM caching.
-    VERSION 0.0.20240624
-    GIT_TAG fb18ee251b81a43fd9a0f45f2f34770389e916d0
+    VERSION 1.5.2
+    GIT_TAG 24d1807a9a3e0df81103a0be9be7ad28ee34c85a
     GIT_REPOSITORY "https://github.com/openxla/stablehlo.git"
     PATCHES "${MLIR_TENSORRT_ROOT_DIR}/build_tools/stablehlo.patch"
     OPTIONS

diff --git a/mlir-tensorrt/README.md b/mlir-tensorrt/README.md
@@ -23,27 +23,66 @@ We currently support only building on Linux x86 systems.
 We support building several different ways (only via CMake) depending on use-case.
 
 In each case, the LLVM-Project version that we are currently aligned to is
-given in `build_tools/llvm_commit.txt`.
+given in `build_tools/cmake/LLVMCommit.txt`.
 
 Note that currently we provide an LLVM patch which essentially cherry-picks the
 bug fixes from [this open MLIR PR](https://github.com/llvm/llvm-project/pull/91524).
 
-1. Build as a Standalone Project with LLVM provided
-2. Build as a Standalone Project with LLVM downloaded by CMake
+1. Build as a Standalone Project with LLVM downloaded by CMake
+2. Build as a Standalone Project with LLVM provided by User
 3. Build as a sub-project of a larger build (e.g. `add_subdirectory`)
 4. Build via LLVM-External-Projects mechanism
 
-Here we only show how to do Option 2.
+Here we only show how to do Option 1 and Option 2.
+
+## Option 1: Build as a Standalone Project with LLVM downloaded by CMake
+
+This is the simplest way to get started and incurs low download overhead
+since we download LLVM-Project as a zip archive directly from GitHub
+at our pinned commit.
+
+```sh
+# Note: we use clang and lld here. These are recommended.
+# However, GNU toolchains will also work,
+# clang toolchain is optional.
+cmake -B ./build -S . -G Ninja \
+    -DCMAKE_BUILD_TYPE=RelWithDebInfo \
+    -DMLIR_TRT_PACKAGE_CACHE_DIR=${PWD}/.cache.cpm \
+    -DMLIR_TRT_ENABLE_ASSERTIONS=ON \
+    -DMLIR_TRT_DOWNLOAD_TENSORRT_VERSION=10.2 \
+    -DCMAKE_C_COMPILER=clang -DCMAKE_CXX_COMPILER=clang++ \
+    -DMLIR_TRT_USE_LINKER=lld
+
+# Example build commands:
+
+# Build everything
+ninja -C build all
+
+# Build and run tests
+ninja -C build/mlir-tensorrt check-mlir-executor
+ninja -C build/mlir-tensorrt check-mlir-tensorrt-dialect
+ninja -C build/mlir-tensorrt check-mlir-tensorrt
+
+# Build wheels (output in `build/wheels`)
+ninja -C build mlir-tensorrt-all-wheels
+```
+
+## Option 2: Build as a Standalone Project with LLVM provided by User
+
+This is more complex but lets you "bring your own LLVM-Project" source
+code or binary.
+
 
 1. Build MLIR
 
 ```sh
 # Clone llvm-project
 git clone https://github.com/llvm/llvm-project.git llvm-project
 
-# Checkout the right commit
+# Checkout the right commit. Of course, you may try
+# a newer commit or your own modified LLVM-Project.
 cd llvm-project
-git checkout $(cat ../build_tools/llvm_commit.txt)
+git checkout $(cat build_tools/cmake/LLVMCommit.cmake | grep -Po '(?<=").*(?=")')
 
 # Apply patch from llvm-project PR 91524
 git apply ../build_tools/llvm-project.patch

diff --git a/mlir-tensorrt/build_tools/cmake/LLVMCommit.cmake b/mlir-tensorrt/build_tools/cmake/LLVMCommit.cmake
@@ -0,0 +1 @@
+set(LLVM_COMMIT "c8b5d30f707757a4fe4d9d0bb01f762665f6942f")
diff --git a/mlir-tensorrt/build_tools/cmake/ManagedLLVM.cmake b/mlir-tensorrt/build_tools/cmake/ManagedLLVM.cmake
@@ -1,16 +1,6 @@
 macro(mtrt_llvm_project)
   CPMAddPackage(
-    NAME llvm_project
-    URL "https://github.com/llvm/llvm-project/archive/${MLIR_TRT_LLVM_COMMIT}.zip"
-    EXCLUDE_FROM_ALL TRUE
-    SOURCE_SUBDIR "llvm"
-    PATCHES "${CMAKE_SOURCE_DIR}/build_tools/llvm-project.patch"
-    OPTIONS
-      "LLVM_ENABLE_PROJECTS mlir"
-      "MLIR_ENABLE_BINDINGS_PYTHON ${MLIR_TRT_ENABLE_PYTHON}"
-      "LLVM_TARGETS_TO_BUILD host"
-      "LLVM_ENABLE_ASSERTIONS ${MLIR_TRT_ENABLE_ASSERTIONS}"
-      "LLVM_USE_LINKER ${MLIR_TRT_USE_LINKER}"
+    ${ARGN}
     )
   set(MLIR_CMAKE_DIR "${CMAKE_BINARY_DIR}/lib/cmake/mlir")
   set(LLVM_CMAKE_DIR "${llvm_project_BINARY_DIR}/lib/cmake/llvm")

diff --git a/mlir-tensorrt/build_tools/llvm_commit.txt b/mlir-tensorrt/build_tools/llvm_commit.txt
diff --git a/mlir-tensorrt/build_tools/scripts/cicd_build.sh b/mlir-tensorrt/build_tools/scripts/cicd_build.sh
@@ -3,7 +3,7 @@ set -ex
 set -o pipefail
 
 REPO_ROOT=$(pwd)
-BUILD_DIR="${BUILD_DIR:=${REPO_ROOT}/build/mlir-tensorrt}"
+BUILD_DIR="${BUILD_DIR:=${REPO_ROOT}/build}"
 
 ENABLE_NCCL=${ENABLE_NCCL:OFF}
 RUN_LONG_TESTS=${RUN_LONG_TESTS:-False}
@@ -23,7 +23,6 @@ cmake -GNinja -B "${BUILD_DIR}" -S "${REPO_ROOT}" \
   -DMLIR_TRT_DOWNLOAD_TENSORRT_VERSION="$DOWNLOAD_TENSORRT_VERSION" \
   -DLLVM_LIT_ARGS="${LLVM_LIT_ARGS}" \
   -DENABLE_ASAN="${ENABLE_ASAN}" \
-  -DMLIR_DIR=${REPO_ROOT}/build/llvm-project/lib/cmake/mlir \
   -DCMAKE_PLATFORM_NO_VERSIONED_SONAME=ON
 
 echo "==== Running Build ==="

diff --git a/mlir-tensorrt/build_tools/stablehlo.patch b/mlir-tensorrt/build_tools/stablehlo.patch
@@ -1,10 +1,34 @@
+diff --git a/stablehlo/integrations/c/CMakeLists.txt b/stablehlo/integrations/c/CMakeLists.txt
+index 014a1a8d..1e0bcfb4 100644
+--- a/stablehlo/integrations/c/CMakeLists.txt
++++ b/stablehlo/integrations/c/CMakeLists.txt
+@@ -28,7 +28,6 @@ add_mlir_public_c_api_library(ChloCAPI
+
+   LINK_LIBS PUBLIC
+   ChloOps
+-  LLVMSupport
+ )
+
+ add_mlir_public_c_api_library(StablehloCAPI
+diff --git a/stablehlo/reference/CMakeLists.txt b/stablehlo/reference/CMakeLists.txt
+index b3a406b0..af6c38da 100644
+--- a/stablehlo/reference/CMakeLists.txt
++++ b/stablehlo/reference/CMakeLists.txt
+@@ -169,7 +169,6 @@ add_mlir_library(StablehloReferenceProcessGrid
+   ProcessGrid.cpp
+
+   LINK_LIBS PUBLIC
+-  LLVMSupport
+   MLIRIR
+   MLIRSupport
+   StablehloReferenceTensor
 diff --git a/stablehlo/transforms/StablehloAggressiveSimplification.cpp b/stablehlo/transforms/StablehloAggressiveSimplification.cpp
-index 8e9332f3..2bfdf34f 100644
+index 6445f72d..cc6c542d 100644
 --- a/stablehlo/transforms/StablehloAggressiveSimplification.cpp
 +++ b/stablehlo/transforms/StablehloAggressiveSimplification.cpp
-@@ -39,6 +39,40 @@ namespace stablehlo {
+@@ -58,6 +58,40 @@ namespace stablehlo {
  #include "stablehlo/transforms/Passes.h.inc"
- 
+
  namespace {
 +
 +/// Does the same as PatternRewriter::replaceOpWithNewOp, but with a twist.
@@ -43,34 +67,34 @@ index 8e9332f3..2bfdf34f 100644
  // This is an upper limit on how many elements canonicalization patterns are
  // allowed to materialize as new constants.
  constexpr int64_t kFoldOpEltLimit = 65536;
-@@ -119,7 +153,7 @@ struct AddOpCanon final : OpRewritePattern<mlir::stablehlo::AddOp> {
+@@ -138,7 +172,7 @@ struct AddOpCanon final : OpRewritePattern<mlir::stablehlo::AddOp> {
      if (TypedAttr res;
          lhsAttr && rhsAttr &&
          (res = foldBinaryOpIntOrFloat(lhsAttr, rhsAttr, std::plus<>{}))) {
 -      rewriter.replaceOpWithNewOp<mlir::stablehlo::ConstantOp>(op, res);
 +      refineOpWithNewOp<mlir::stablehlo::ConstantOp>(rewriter, op, res);
        return success();
      }
- 
-@@ -157,7 +191,7 @@ struct SubtractOpCanon final : OpRewritePattern<mlir::stablehlo::SubtractOp> {
+
+@@ -176,7 +210,7 @@ struct SubtractOpCanon final : OpRewritePattern<mlir::stablehlo::SubtractOp> {
      if (TypedAttr res;
          lhsAttr && rhsAttr &&
          (res = foldBinaryOpIntOrFloat(lhsAttr, rhsAttr, std::minus<>{}))) {
 -      rewriter.replaceOpWithNewOp<mlir::stablehlo::ConstantOp>(op, res);
 +      refineOpWithNewOp<mlir::stablehlo::ConstantOp>(rewriter, op, res);
        return success();
      }
- 
-@@ -208,7 +242,7 @@ struct MulOpCanon final : OpRewritePattern<mlir::stablehlo::MulOp> {
+
+@@ -227,7 +261,7 @@ struct MulOpCanon final : OpRewritePattern<mlir::stablehlo::MulOp> {
      if (TypedAttr res;
          lhsAttr && rhsAttr &&
          (res = foldBinaryOpIntOrFloat(lhsAttr, rhsAttr, std::multiplies<>{}))) {
 -      rewriter.replaceOpWithNewOp<mlir::stablehlo::ConstantOp>(op, res);
 +      refineOpWithNewOp<mlir::stablehlo::ConstantOp>(rewriter, op, res);
        return success();
      }
- 
-@@ -299,15 +333,16 @@ struct CompareOpCanon final : OpRewritePattern<mlir::stablehlo::CompareOp> {
+
+@@ -318,15 +352,16 @@ struct CompareOpCanon final : OpRewritePattern<mlir::stablehlo::CompareOp> {
          case ComparisonDirection::EQ:
          case ComparisonDirection::GE:
          case ComparisonDirection::LE: {
@@ -91,49 +115,27 @@ index 8e9332f3..2bfdf34f 100644
            return success();
          }
        }
-@@ -336,7 +371,7 @@ struct CompareOpCanon final : OpRewritePattern<mlir::stablehlo::CompareOp> {
+@@ -355,7 +390,7 @@ struct CompareOpCanon final : OpRewritePattern<mlir::stablehlo::CompareOp> {
               [direction, kind = *compType](const APInt &a, const APInt &b) {
                 return calculateComp(kind, direction, a, b);
               }))) {
 -      rewriter.replaceOpWithNewOp<mlir::stablehlo::ConstantOp>(op, res);
 +      refineOpWithNewOp<mlir::stablehlo::ConstantOp>(rewriter, op, res);
        return success();
      }
-
-@@ -349,8 +384,6 @@ struct SelectOpCanon final : OpRewritePattern<mlir::stablehlo::SelectOp> {
-
-   LogicalResult matchAndRewrite(mlir::stablehlo::SelectOp op,
-                                 PatternRewriter &rewriter) const override {
--    RankedTensorType type = op.getType();
--
-     Value trueVal = op.getOnTrue();
-     Value falseVal = op.getOnFalse();
-
-@@ -373,7 +406,9 @@ struct SelectOpCanon final : OpRewritePattern<mlir::stablehlo::SelectOp> {
-
-     // Handle elementwise selection when both outcomes are also constants. This
-     // will create a new, likely non-splat constant.
--    if (cond.getNumElements() > kFoldOpEltLimit) return failure();
-+    if (!cond.getShapedType().hasStaticShape() ||
-+        cond.getNumElements() > kFoldOpEltLimit)
-+      return failure();
-
-     ElementsAttr trueAttr;
-     if (!matchPattern(trueVal, m_Constant(&trueAttr))) return failure();
-@@ -389,8 +424,10 @@ struct SelectOpCanon final : OpRewritePattern<mlir::stablehlo::SelectOp> {
+
+@@ -408,8 +443,8 @@ struct SelectOpCanon final : OpRewritePattern<mlir::stablehlo::SelectOp> {
        newValues.push_back(condElem ? trueElem : falseElem);
      }
- 
+
 -    rewriter.replaceOpWithNewOp<mlir::stablehlo::ConstantOp>(
 -        op, DenseElementsAttr::get(type, newValues));
 +    refineOpWithNewOp<mlir::stablehlo::ConstantOp>(
-+        rewriter, op,
-+        DenseElementsAttr::get(cast<RankedTensorType>(trueVal.getType()),
-+                               newValues));
++        rewriter, op, DenseElementsAttr::get(type, newValues));
      return success();
    }
  };
-@@ -465,9 +502,10 @@ struct BroadcastInDimOpCanon final
+@@ -484,9 +519,10 @@ struct BroadcastInDimOpCanon final
      // Handle splat broadcasts.
      if (SplatElementsAttr cstAttr;
          matchPattern(operand, m_Constant(&cstAttr))) {
@@ -146,10 +148,10 @@ index 8e9332f3..2bfdf34f 100644
 +                                 cstAttr.getSplatValue<Attribute>()));
        return success();
      }
- 
-@@ -538,8 +576,8 @@ struct ConcatenateOpCanon final
+
+@@ -557,8 +593,8 @@ struct ConcatenateOpCanon final
      }
- 
+
      assert(newElems.size() == numElems);
 -    rewriter.replaceOpWithNewOp<mlir::stablehlo::ConstantOp>(
 -        op, DenseElementsAttr::get(op.getType(), newElems));
@@ -158,10 +160,10 @@ index 8e9332f3..2bfdf34f 100644
      return success();
    }
  };
-@@ -557,39 +595,6 @@ struct ConvertOpCanon final : OpRewritePattern<mlir::stablehlo::ConvertOp> {
+@@ -576,39 +612,6 @@ struct ConvertOpCanon final : OpRewritePattern<mlir::stablehlo::ConvertOp> {
    }
  };
- 
+
 -/// Does the same as PatternRewriter::replaceOpWithNewOp, but with a twist.
 -///
 -/// Sometimes, we want to replace an op with a new op and simultaneously refine
@@ -198,8 +200,8 @@ index 8e9332f3..2bfdf34f 100644
  /// If a DynamicBroadCastInDimOp is not actually dynamic, use an ordinary
  /// BroadcastInDimOp.
  struct DynamicBroadcastInDimOpNotActuallyDynamic final
-@@ -924,8 +929,8 @@ struct GetDimensionSizeOpCanon final
- 
+@@ -943,8 +946,8 @@ struct GetDimensionSizeOpCanon final
+
      auto elemTy = cast<IntegerType>(op.getType().getElementType());
      IntegerAttr elemVal = rewriter.getIntegerAttr(elemTy, dimSize);
 -    rewriter.replaceOpWithNewOp<mlir::stablehlo::ConstantOp>(
@@ -209,9 +211,9 @@ index 8e9332f3..2bfdf34f 100644
      return success();
    }
  };
-@@ -1017,16 +1022,17 @@ struct ReshapeOpCanon final : OpRewritePattern<mlir::stablehlo::ReshapeOp> {
+@@ -1036,16 +1039,17 @@ struct ReshapeOpCanon final : OpRewritePattern<mlir::stablehlo::ReshapeOp> {
      if (!matchPattern(op.getOperand(), m_Constant(&cstAttr))) return failure();
- 
+
      if (auto splat = dyn_cast<SplatElementsAttr>(cstAttr)) {
 -      rewriter.replaceOpWithNewOp<mlir::stablehlo::ConstantOp>(
 -          op, SplatElementsAttr::get(op.getType(),
@@ -222,7 +224,7 @@ index 8e9332f3..2bfdf34f 100644
 +                                 splat.getSplatValue<Attribute>()));
        return success();
      }
- 
+
      auto elements =
          llvm::to_vector_of<Attribute>(cstAttr.getValues<Attribute>());
 -    rewriter.replaceOpWithNewOp<mlir::stablehlo::ConstantOp>(
@@ -232,3 +234,15 @@ index 8e9332f3..2bfdf34f 100644
      return success();
    }
  };
+diff --git a/stablehlo/transforms/StablehloRefineShapes.cpp b/stablehlo/transforms/StablehloRefineShapes.cpp
+index e5120133..7118a9e7 100644
+--- a/stablehlo/transforms/StablehloRefineShapes.cpp
++++ b/stablehlo/transforms/StablehloRefineShapes.cpp
+@@ -151,7 +151,7 @@ LogicalResult refineReturnTypes(PatternRewriter& rewriter, Operation* op,
+   // upstream API to achieve this directly, but if it's introduced in the
+   // future, we could use it here.
+   rewriter.replaceOpUsesWithIf(op, op->getResults(),
+-                               [](OpOperand& use) { return false; });
++                               [](OpOperand& use) { return true; });
+   return success();
+ }
Original file line number	Diff line number	Diff line change
		@@ -0,0 +1 @@
		set(LLVM_COMMIT "c8b5d30f707757a4fe4d9d0bb01f762665f6942f")