Integrate internal changes

mlir-tensorrt/CMakeLists.txt

@@ -190,6 +190,8 @@ if(MLIR_TRT_ENABLE_PYTHON)
   mlir_tensorrt_find_dlpack()
 endif()
 
+find_package(MLIRTensorRTCommon REQUIRED)
+
 #--------------------------------------------------
 # Diagnostics
 #--------------------------------------------------

mlir-tensorrt/DependencyProvider.cmake

@@ -17,7 +17,7 @@ if("${MLIR_TRT_USE_LLVM}" STREQUAL "prebuilt")
   set(MTRT_BUILD_LLVM_FROM_SOURCE OFF)
 endif()
 
-set(MLIR_TRT_LLVM_COMMIT "729416e586fba71b4f63d71b1b5c765aefbf200b")
+set(MLIR_TRT_LLVM_COMMIT "f137c3d592e96330e450a8fd63ef7e8877fc1908")
 
 set(mlir_patch_dir "${CMAKE_CURRENT_LIST_DIR}/build_tools/patches/mlir")
 
@@ -43,7 +43,7 @@ else()
       "${mlir_patch_dir}/0005-mlir-memref-Fix-memref.global-overly-constrained-ver.patch"
       "${mlir_patch_dir}/0006-mlir-emitc-Fix-two-EmitC-bugs.patch"
       "${mlir_patch_dir}/0009-mlir-Support-FileLineColRange-in-LLVM-debug-translat.patch"
-      "${mlir_patch_dir}/0010-MLIR-Fix-LLVMIRTransforms-build-failure-125485.patch"
+      "${mlir_patch_dir}/0011-MLIR-Fix-bufferization-interface-for-tensor-reshape.patch"
     # Set the CPM cache key to the Git hash for easy navigation.
     PRE_ADD_HOOK [[
       list(APPEND _vap_UNPARSED_ARGUMENTS
@@ -63,6 +63,8 @@ else()
       )
       list(APPEND CMAKE_MODULE_PATH "${MLIR_CMAKE_DIR}")
 
+      set(MLIR_MAIN_SRC_DIR "${LLVM_SOURCE_DIR}/mlir" CACHE STRING "" FORCE)
+
       if(TARGET MLIRPythonExtension.Core)
         get_property(mlir_core_pybind_capi_embed
           TARGET MLIRPythonExtension.Core
@@ -102,14 +104,12 @@ set(stablehlo_patch_dir "${CMAKE_SOURCE_DIR}/build_tools/patches/stablehlo")
 nv_register_package(
   NAME Stablehlo
   VERSION 1.9.3
-  GIT_TAG 459897561d365ef97caba46984847f9184d472ec
+  GIT_TAG 4bf77d23bd9150782a70d85fda9c12a2dec5328c
   GIT_REPOSITORY "https://github.com/openxla/stablehlo.git"
   PATCHES
     "${stablehlo_patch_dir}/0001-Fix-a-couple-missing-checks-for-static-shapes-in-sta.patch"
     "${stablehlo_patch_dir}/0002-cmake-Update-usage-of-HandleLLVMOptions-and-LLVM_DEF.patch"
-    "${stablehlo_patch_dir}/0003-Don-t-insert-unnecessary-arith.index_cast-ops.patch"
     "${stablehlo_patch_dir}/0004-Fix-ZeroExtent-condition-in-simplification-pattern.patch"
-    "${stablehlo_patch_dir}/0005-Fix-crash-on-ComplexType-in-PointwiseToLinalgMapConv.patch"
     "${stablehlo_patch_dir}/0006-Remove-explicit-use-of-LLVMSupport.patch"
     "${stablehlo_patch_dir}/0007-Fix-circular-dependence-between-StablehloPasses-and-.patch"
   OPTIONS
@@ -123,6 +123,42 @@ nv_register_package(
   ]]
 )
 
+#-------------------------------------------------------------------------------------
+# MLIRTensorRTCommon
+#
+# MLIRTensorRTCommon is a sub-project that contains components used across the
+# other sub-projects like MLIRExecutor and MLIRTensorRTDialect.
+#-------------------------------------------------------------------------------------
+
+nv_register_package(
+  NAME MLIRTensorRTCommon
+  SOURCE_DIR "${CMAKE_SOURCE_DIR}/common"
+)
+
+# -----------------------------------------------------------------------------
+# NVTX
+# -----------------------------------------------------------------------------
+
+nv_register_package(
+  NAME NVTX
+  GIT_REPOSITORY https://github.com/NVIDIA/NVTX.git
+  GIT_TAG v3.1.0
+  GIT_SHALLOW TRUE
+  SOURCE_SUBDIR c
+  EXCLUDE_FROM_ALL TRUE
+  DOWNLOAD_ONLY TRUE
+  POST_ADD_HOOK [[
+    if(NOT TARGET nvtx3-cpp)
+      add_library(nvtx3-cpp INTERFACE IMPORTED)
+      target_include_directories(nvtx3-cpp INTERFACE
+        "$<BUILD_INTERFACE:${NVTX_SOURCE_DIR}/c/include>")
+      # Ignore some warnings due to NVTX3 code style.
+      target_compile_options(nvtx3-cpp INTERFACE
+        -Wno-missing-braces)
+    endif()
+  ]]
+)
+
 #-------------------------------------------------------------------------------------
 # MLIR-Executor
 #
@@ -158,26 +194,23 @@ nv_register_package(
 #-------------------------------------------------------------------------------------
 # Torch-MLIR
 #-------------------------------------------------------------------------------------
-set(torch_mlir_patch_dir "${CMAKE_SOURCE_DIR}/build_tools/patches/torch_mlir")
-
+set(MLIR_TRT_TORCH_MLIR_COMMIT "9f2ba5abaa85cefd95cc85579fafd0c53c1101e8")
 nv_register_package(
   NAME torch_mlir
-  GIT_REPOSITORY https://github.com/llvm/torch-mlir.git
-  GIT_TAG 0bb263e99415d43255350d29263097b4980303bf
-  PATCHES 
-    "build_tools/patches/torch_mlir/0001-cmake-Allow-finding-Stablehlo-via-find_package.patch"
-    "build_tools/patches/torch_mlir/0002-Make-compatible-with-more-recent-Stablehlo-version.patch"
-    "build_tools/patches/torch_mlir/0003-Fix-some-configuration-paths-in-LIT-cfg.patch"
-  EXCLUDE_FROM_ALL TRUE
+  URL "https://github.com/llvm/torch-mlir/archive/${MLIR_TRT_TORCH_MLIR_COMMIT}.zip"
   # We need to specify an existing directory that is not actually a submodule
   # since GIT_SUBMODULES does not except the empty string due to
   # https://gitlab.kitware.com/cmake/cmake/-/issues/24578
   GIT_SUBMODULES "docs"
-  OPTIONS
-    "TORCH_MLIR_OUT_OF_TREE_BUILD ON"
-    "TORCH_MLIR_ENABLE_STABLEHLO ON"
-    "TORCH_MLIR_EXTERNAL_STABLEHLO_DIR find_package"
-    "TORCH_MLIR_ENABLE_TOSA OFF"
+  DOWNLOAD_ONLY TRUE
+
+  POST_ADD_HOOK [[
+    add_subdirectory(
+        ${CMAKE_SOURCE_DIR}/third_party/torch-mlir-cmake
+        ${CMAKE_BINARY_DIR}/_deps/torch_mlir-build
+        EXCLUDE_FROM_ALL
+    )
+  ]]
 )
 
 #-------------------------------------------------------------------------------------
@@ -202,7 +235,7 @@ macro(mtrt_provide_dependency method dep_name)
   endif()
 
   if("${dep_name}" MATCHES
-     "^(MLIRExecutor|MLIRTensorRTDialect|Stablehlo|torch_mlir)$")
+     "^(MLIRExecutor|MLIRTensorRTDialect|Stablehlo|torch_mlir|NVTX|MLIRTensorRTCommon)$")
     nv_add_package("${dep_name}")
     set("${dep_name}_FOUND" TRUE)
   endif()
@@ -230,6 +263,7 @@ macro(mtrt_provide_dependency method dep_name)
       find_package(LLVM ${ARGN} BYPASS_PROVIDER)
     endif()
   endif()
+
 endmacro()
 
 cmake_language(

mlir-tensorrt/build_tools/cmake/Dependencies.cmake

@@ -6,10 +6,16 @@ include(${CMAKE_CURRENT_LIST_DIR}/TensorRTDownloadURL.cmake)
 # expected version.
 #-------------------------------------------------------------------------------------
 macro(get_tensorrt_version nvinfer_version_file out_var)
-  file(STRINGS "${nvinfer_version_file}" VERSION_STRINGS REGEX "#define NV_TENSORRT_.*")
+  file(STRINGS "${nvinfer_version_file}" VERSION_STRINGS REGEX "#define (TRT_.+|NV_TENSORRT_.+) [0-9]+")
   foreach(TYPE MAJOR MINOR PATCH BUILD)
-    string(REGEX MATCH "NV_TENSORRT_${TYPE} [0-9]+" TRT_TYPE_STRING ${VERSION_STRINGS})
-    string(REGEX MATCH "[0-9]+" TRT_${TYPE} ${TRT_TYPE_STRING})
+    string(REGEX MATCH "(TRT_${TYPE}_ENTERPRISE|NV_TENSORRT_${TYPE}) [0-9]+" TRT_TYPE_STRING ${VERSION_STRINGS})
+    if("${TRT_TYPE_STRING}" STREQUAL "")
+      message(FATAL_ERROR "Failed to extract TensorRT ${TYPE} version from ${nvinfer_version_file}")
+    endif()
+    string(REGEX MATCH "[0-9]+" "TRT_${TYPE}" "${TRT_TYPE_STRING}")
+    if("TRT_${TYPE}" STREQUAL "")
+      message(FATAL_ERROR "Failed to extract TensorRT ${TYPE} version from ${nvinfer_version_file}")
+    endif()
   endforeach(TYPE)
   set("${out_var}" "${TRT_MAJOR}.${TRT_MINOR}.${TRT_PATCH}.${TRT_BUILD}")
 endmacro()
@@ -50,7 +56,7 @@ macro(configure_tensorrt_python_plugin_header)
   if(ARG_INSTALL_DIR)
     find_file(
       trt_python_plugin_header
-      NAMES plugin.h
+      NAMES NvInferPythonPlugin.h plugin.h
       HINTS ${ARG_INSTALL_DIR} ${ARG_INSTALL_DIR}/python/include/impl
       PATHS ${ARG_INSTALL_DIR} ${ARG_INSTALL_DIR}/python/include/impl
       REQUIRED
@@ -60,7 +66,7 @@ macro(configure_tensorrt_python_plugin_header)
   else()
     find_path(
       trt_python_plugin_header
-      NAMES plugin.h
+      NAMES NvInferPythonPlugin.h plugin.h
       REQUIRED
       NO_CACHE
     )
@@ -173,36 +179,6 @@ function(find_tensorrt)
     )
 endfunction()
 
-macro(configure_tensorrt_python_plugin_header)
-  if(ARG_INSTALL_DIR)
-    find_file(
-      trt_python_plugin_header
-      NAMES plugin.h
-      HINTS ${ARG_INSTALL_DIR} ${ARG_INSTALL_DIR}/python/include/impl
-      PATHS ${ARG_INSTALL_DIR} ${ARG_INSTALL_DIR}/python/include/impl
-      REQUIRED
-      NO_CMAKE_PATH NO_DEFAULT_PATH
-      NO_CACHE
-    )
-  else()
-    find_path(
-      trt_python_plugin_header
-      NAMES plugin.h
-      REQUIRED
-      NO_CACHE
-    )
-  endif()
-  file(MAKE_DIRECTORY "${CMAKE_BINARY_DIR}/include/nvinfer")
-  file(COPY_FILE "${trt_python_plugin_header}"
-    "${CMAKE_BINARY_DIR}/include/nvinfer/trt_plugin_python.h"
-    ONLY_IF_DIFFERENT
-    RESULT copy_result
-  )
-  if(copy_result)
-    message(FATAL_ERROR "failed to copy TensorRT QDP plugin header: ${copy_result}")
-  endif()
-endmacro()
-
 #-------------------------------------------------------------------------------------
 # Download and add DLPack to the build (header only)
 #-------------------------------------------------------------------------------------

mlir-tensorrt/build_tools/patches/mlir/0005-mlir-memref-Fix-memref.global-overly-constrained-ver.patch

@@ -1,18 +1,17 @@
-From f014186374bb3e71d44648781dc03aaefd29f0d5 Mon Sep 17 00:00:00 2001
-From: Christopher Bate <[email protected]>
-Date: Fri, 10 May 2024 22:39:44 -0600
-Subject: [PATCH 05/10] [mlir][memref] Fix memref.global overly constrained
- verifier check
+From 07f534dac7f915a496265c14745c0bc643185efe Mon Sep 17 00:00:00 2001
+From: Sagar Shelke <[email protected]>
+Date: Tue, 1 Jul 2025 00:17:04 +0000
+Subject: [PATCH] Apply patch 0005
 
 ---
  mlir/lib/Dialect/MemRef/IR/MemRefOps.cpp | 10 ++++++----
  1 file changed, 6 insertions(+), 4 deletions(-)
 
 diff --git a/mlir/lib/Dialect/MemRef/IR/MemRefOps.cpp b/mlir/lib/Dialect/MemRef/IR/MemRefOps.cpp
-index 4f75b7618d63..f12f41437759 100644
+index 11597505e788..66ce4b3638b0 100644
 --- a/mlir/lib/Dialect/MemRef/IR/MemRefOps.cpp
 +++ b/mlir/lib/Dialect/MemRef/IR/MemRefOps.cpp
-@@ -1117,7 +1117,7 @@ struct DimOfMemRefReshape : public OpRewritePattern<DimOp> {
+@@ -1124,7 +1124,7 @@ struct DimOfMemRefReshape : public OpRewritePattern<DimOp> {
          }
        } // else dim.getIndex is a block argument to reshape->getBlock and
          // dominates reshape
@@ -21,7 +20,7 @@ index 4f75b7618d63..f12f41437759 100644
      else if (dim->getBlock() != reshape->getBlock() &&
               !dim.getIndex().getParentRegion()->isProperAncestor(
                   reshape->getParentRegion())) {
-@@ -1607,9 +1607,11 @@ LogicalResult GlobalOp::verify() {
+@@ -1614,9 +1614,11 @@ LogicalResult GlobalOp::verify() {
      // Check that the type of the initial value is compatible with the type of
      // the global variable.
      if (auto elementsAttr = llvm::dyn_cast<ElementsAttr>(initValue)) {
@@ -37,5 +36,5 @@ index 4f75b7618d63..f12f41437759 100644
                 << tensorType << ", but was of type " << initType;
      }
 -- 
-2.46.0
+2.48.1
 

mlir-tensorrt/build_tools/patches/mlir/0006-mlir-emitc-Fix-two-EmitC-bugs.patch

@@ -1,7 +1,7 @@
-From 47c84211f72fb407d72e2c8f87019802cda30432 Mon Sep 17 00:00:00 2001
-From: Christopher Bate <cbate@nvidia.com>
-Date: Mon, 27 Jan 2025 08:28:33 +0000
-Subject: [PATCH 06/10] [mlir][emitc] Fix two EmitC bugs
+From d81aaed8cb0190807dbe378e469fde53101f32eb Mon Sep 17 00:00:00 2001
+From: Sagar Shelke <shelkesagar29@yahoo.com>
+Date: Tue, 1 Jul 2025 00:19:03 +0000
+Subject: [PATCH] Apply patch 0006
 
 ---
  .../mlir/Conversion/FuncToEmitC/FuncToEmitC.h |  4 +-
@@ -112,18 +112,18 @@ index 0b97f2641ad0..d2f368a7148d 100644
    if (failed(
            applyPartialConversion(getOperation(), target, std::move(patterns))))
 diff --git a/mlir/lib/Target/Cpp/TranslateToCpp.cpp b/mlir/lib/Target/Cpp/TranslateToCpp.cpp
-index 01de0e41f203..4f600f92ba6d 100644
+index b00820ffc542..803c58cc35c6 100644
 --- a/mlir/lib/Target/Cpp/TranslateToCpp.cpp
 +++ b/mlir/lib/Target/Cpp/TranslateToCpp.cpp
-@@ -273,6 +273,7 @@ private:
+@@ -282,6 +282,7 @@ private:
    ExpressionOp emittedExpression;
    SmallVector<int> emittedExpressionPrecedence;
 
 +public:
    void pushExpressionPrecedence(int precedence) {
      emittedExpressionPrecedence.push_back(precedence);
    }
-@@ -670,12 +671,14 @@ static LogicalResult printOperation(CppEmitter &emitter,
+@@ -695,12 +696,14 @@ static LogicalResult printOperation(CppEmitter &emitter,
      if (auto t = dyn_cast<IntegerAttr>(attr)) {
        // Index attributes are treated specially as operand index.
        if (t.getType().isIndex()) {
@@ -143,5 +143,5 @@ index 01de0e41f203..4f600f92ba6d 100644
        }
      }
 -- 
-2.46.0
+2.48.1
 

mlir-tensorrt/build_tools/patches/mlir/0009-mlir-Support-FileLineColRange-in-LLVM-debug-translat.patch

@@ -1,15 +1,14 @@
-From 51c99ccf1a291295aed12a36395760026c268cbb Mon Sep 17 00:00:00 2001
-From: Christopher Bate <[email protected]>
-Date: Tue, 11 Mar 2025 22:34:24 +0000
-Subject: [PATCH 09/10] [mlir] Support FileLineColRange in LLVM debug
- translation
+From d679ee8c7978fe63321e90c5c6583b604ca3d1a5 Mon Sep 17 00:00:00 2001
+From: Sagar Shelke <[email protected]>
+Date: Tue, 1 Jul 2025 00:25:43 +0000
+Subject: [PATCH] Apply patch 0009
 
 ---
  mlir/lib/Target/LLVMIR/DebugTranslation.cpp | 8 ++++++++
  1 file changed, 8 insertions(+)
 
 diff --git a/mlir/lib/Target/LLVMIR/DebugTranslation.cpp b/mlir/lib/Target/LLVMIR/DebugTranslation.cpp
-index cf734de49acd..c55d9a204468 100644
+index 1d3ed6f3262f..93e1d08faf4f 100644
 --- a/mlir/lib/Target/LLVMIR/DebugTranslation.cpp
 +++ b/mlir/lib/Target/LLVMIR/DebugTranslation.cpp
 @@ -547,6 +547,14 @@ llvm::DILocation *DebugTranslation::translateLoc(Location loc,
@@ -28,5 +27,5 @@ index cf734de49acd..c55d9a204468 100644
      ArrayRef<Location> locations = fusedLoc.getLocations();
 
 -- 
-2.46.0
+2.48.1