Integrate internal changes; LLVM upgrade (#567)

This change integrates the following internal changes (listed from
oldest to newest):

--

ce06fa5b940804009ebfb2652552241057318485 by Samurdhi Karunaratne
<[email protected]>

Support for TRT AOT-Quickly Deployable Plugins (AOT-QDPs)

Adds support for TRT's AOT QDPs, which are available from TRT 10.9
onwards. Guards have been added so that code requiring TRT >= 10.9 won't
break builds against earlier TRT releases.

The changes mostly focus on adding if-else branches for a few
specialized methods in AOT-QDP plugin and plugin creators classes.

--
3964df047bc7ab266054701458709a16c87d7d3d by Chris Bate
<[email protected]>:

NFC: Fixup Executor CMakeLists

Remove outdated LLVM configuration logic

--
a55a19e6cd91f326d4d52372bbcfb89e646b3ee9 by Sagar Shelke
<[email protected]>:

[Dialect/Plan] Add `i16` as valid type after clustering

This MR adds `i16` as valid type after clustering in
`PostClusteringValidation`
pass. Previously op and op type both were checked in single if
condition thus it was hard to tell whether op is problematic or op type
is.
Now, those checks are separated with clear error message.

--
a3b30f5c399984d97903c7baa4952554a4e29eed by Sagar Shelke
<[email protected]>:

[Conversion/StablehloToTensorRT] Fix a bug in `numBatchDims` calculation
in `stablehlo.reduce_window` conversion

This MR fixes an issue in `stablehlo.reduce_window` converter.
Previously, all unit dimensions in window dimensions were considered
batch dimensions (i.e. non spatial dims, strictly speaking). This is
incorrect, since it is possible for spatial dimension to be unit. Based
on TensorRT's requirements, batch dimensions is set to `rank(input) -
spatial_dims`,
where spatial dims is 2 for 4D input and 3 for 5D input.

Valid MLIR test is added.

--
eb38f32f804db10b61f619c0da81a331f83b9c1e by Chris Bate
<[email protected]>:

[compiler][StableHloExt] Fix gather canonicalization in the presence of
batching dims

In Stablehlo v1.1.0, the notion of "batching dimensions" was added to
the
'stablehlo.gather' operation. We did not account for batching dimensions
in
the `stablehlo-ext-canonicalize-gather` pass, which could cause a crash.
This change adds a check so that the patterns in that pass correctly
reject
the gather operations with batching dimensions.

--
90d5be62c8cb36cf80001cc83f9d8e21679ec137 by Chris Bate
<[email protected]>:

[compiler][plan] Fix bug in plan-eliminate-shape-ops pass

Fixes a couple subtle issues in 'plan-eliminate-shape-ops':

- `tensorrt.call` can't have arguments erased "in-place" since it has
  multiple variadic operand segments.
- Procedure for calculating symbol user map had small bug.

--
9c48a5251628ba43067870b4dee4e99477d69f8e by Sagar Shelke
<[email protected]>:

[compiler/lib] Handle 1D `chlo.top_k` conversion

This MR updates the `ChloToTensorRT` conversion to handle 1D
`chlo.top_k`
conversion to `tensorrt.top_k`by rank expanding 1D input. TensorRT TopK
op doesn't support 1D input.
A valid MLIR test case is added.

--
db1cc9d58a4925718099f4e223ba3cdbb2a80b0f by Sagar Shelke
<[email protected]>:

[tensorrt/lib/Target] Handle blob dense resource conversion to TensorRT
weights

This MR adds support to convert `DenseResourceElementsAttr` attribute
with data stored in dialect resource map (within IR) to TensorRT
nvinfer1::Weights.
We did not face this issue before since JAX represents these attributes
in
`elided` form but stablehlo IR coming from Torch-MLIR has dialect
resources.
INT4 case is handled specially considering TensorRT packs two INT4s into
one
INT8. For all other types, weights are copied from the dialect resource.

In function `isValidTensorRTInputType`, i64 is added as valid
TensorRT input type.

MLIR test case is added.

--

6b28981e07a6c0e49275824a39cfbe8e680cf052 by Chris Bate
<[email protected]>:

Integrate LLVM/MLIR upstream changes

The last 10 changes in each repo are shown below. Commits prefixed by
(*) are patches maintained locally. Commits prefixed by (cp) are patches
coresponding to cherry-picks from upstream onto our current pinned
commit/branch.

llvm-project/
- (cp) 75f0d527fe5d Mon Feb 3 11:51:42 2025 +0200 [MLIR] Remove
unnecessary include from MathToEmitC.h to fix build issue (#125466)
- (*) 9149d7f145b8 Mon Mar 10 22:03:05 2025 +0000 [mlir][linalg] don't
rewrite DPS init operands in `linalg-fold-unit-extent-dims` if they are
block arguments
- (*) 47c84211f72f Mon Jan 27 08:28:33 2025 +0000 [mlir][emitc] Fix two
EmitC bugs
- (*) f014186374bb Fri May 10 22:39:44 2024 -0600 [mlir][memref] Fix
memref.global overly constrained verifier check
- ec66c4af0926 Fri Jan 24 10:44:00 2025 -0500 [AMDGPU][True16][CodeGen]
true16 codegen pattern for f16 canonicalize (#122000)
- a94226f9e6f5 Fri Jan 24 10:30:10 2025 -0500 [llvm-ml] Remove unsafe
getCurrentSegmentOnly() call (#123355)

stablehlo/
- (*) 97c312a1 Mon Mar 10 22:51:38 2025 +0000 Fix ZeroExtent condition
in simplification pattern
- (*) d31c6ef5 Mon Mar 10 22:04:05 2025 +0000 Don't insert unnecessary
`arith.index_cast` ops
- (*) f45c6b58 Sat Feb 15 22:02:17 2025 +0000 [cmake] Update usage of
`HandleLLVMOptions` and `LLVM_DEFINITIONS`
- (*) cc45fc70 Wed Nov 27 00:10:11 2024 +0000 Fix a couple missing
checks for static shapes in `stablehlo-aggressive-folder`
- 48a1e14e Tue Jan 28 20:32:38 2025 -0800 Integrate LLVM at
llvm/llvm-project@aa65f93b71de (#2701)
- 7c50d4ef Tue Jan 28 19:47:38 2025 -0600 Add ResultAccuracy to ExpOp
(#2694)
- 8993ef70 Tue Jan 28 15:07:17 2025 -0800 Bump patch version after
integrate 1.8.11 -> 1.8.12 (#2700)
- 37750d07 Tue Jan 28 18:26:15 2025 +0100 Add stablehlo-coreml to
awesome.md (#2699)
- 540d48f2 Mon Jan 27 16:49:56 2025 +0100 Added GoMLX and gopjrt for
awesome.md (#2696)
- b9d4c702 Mon Jan 27 07:49:38 2025 -0800 Add IREE project link to
awesome.md. (#2698)

torch-mlir/
- (*) c220abdb Fri Feb 14 00:39:36 2025 +0000 [cmake] Allow finding
Stablehlo via 'find_package'
- ca6cf285 Fri Jan 31 10:48:06 2025 +0530 Bump llvm to
llvm/llvm-project@5d6d982 (#3994)
- 12250739 Tue Jan 28 10:33:51 2025 +0530 Integrate LLVM at
e2402615a5a76d46a433dfcc1de10b38a1263c9d (#3982)
- af8514c9 Mon Jan 27 09:57:07 2025 -0800 Fix logical `or`/`and` usage
for MSVC compilation. (#3984)
- 4f0b79c5 Fri Jan 24 19:40:46 2025 +0530 build: manually update PyTorch
version (#3977)
- 2564d7af Wed Jan 22 15:59:27 2025 -0800 Add center_point_box=1 support
in NonMaxSuppression. (#3976)
- 481da8d2 Wed Jan 22 12:41:18 2025 -0500 [TOSA] : Fix float to integer
cast for `torch.ops.aten.to` lowering. (#3946)
- 9dd94fbb Wed Jan 22 17:42:25 2025 +0100 Attempt to fix the Python
wheels for Windows (#3979)
- 2fb7d6e4 Tue Jan 21 17:09:17 2025 +0100 Update default Python versions
(#3970)
- 2cc31d6a Tue Jan 21 10:12:39 2025 +0530 Backend-legal-ops argument for
fx lowering (#3956)

--
ef9acb4debd15e407fb70e5aa4440b2af14244c4 by Tom Fogal
<[email protected]>:

[executor] standalone cpp runtime: run the network during enqueue.

The 'mtrt::tensorrt_enqueue' function was missing the actual call to
'IExecutionContext::enqueueV3'.

GitOrigin-RevId: ea86bc480ce59e393c534adae77250cc3ce7d1e6

Author: christopherbate

mlir-tensorrt/DependencyProvider.cmake

@@ -17,15 +17,16 @@ if("${MLIR_TRT_USE_LLVM}" STREQUAL "prebuilt")
   set(MTRT_BUILD_LLVM_FROM_SOURCE OFF)
 endif()
 
-set(MLIR_TRT_LLVM_COMMIT "6c64c8a6f3f77c30745c751d4163ff6bf2fc323b")
+set(MLIR_TRT_LLVM_COMMIT "ec66c4af09263e68d800971906e60afc27d54a06")
+
+set(mlir_patch_dir "${CMAKE_CURRENT_LIST_DIR}/build_tools/patches/mlir")
 
 if(NOT MTRT_BUILD_LLVM_FROM_SOURCE)
   message(WARNING "Using 'find_package' to locate pre-built LLVM. Please set MLIR_DIR to the directory containing MLIRConfig.cmake")
 else()
-  set(patch_dir "${CMAKE_CURRENT_LIST_DIR}/build_tools/patches/mlir")
+
   nv_register_package(
     NAME LLVM
-    VERSION 0.0.20241126
     URL "https://github.com/llvm/llvm-project/archive/${MLIR_TRT_LLVM_COMMIT}.zip"
     EXCLUDE_FROM_ALL TRUE
     SOURCE_SUBDIR "llvm"
@@ -39,23 +40,11 @@ else()
       # Don't mixup LLVM targets with our project's installation/packaging.
       "LLVM_INSTALL_TOOLCHAIN_ONLY ON"
     PATCHES
-      "${patch_dir}/000_fix_bufferization_tensor_encoding_memory_spaces.patch"
-      "${patch_dir}/001-mlir-Add-a-null-pointer-check-in-symbol-lookup-11516.patch"
-      "${patch_dir}/0003-mlir-EmitC-memref-to-emitc-insert-conversion_casts-1.patch"
-      "${patch_dir}/0004-NFC-mlir-emitc-fix-misspelling-in-description-of-emi.patch"
-      "${patch_dir}/0005-emitc-func-Set-default-dialect-to-emitc-116297.patch"
-      "${patch_dir}/0006-MLIR-EmitC-arith-to-emitc-Fix-lowering-of-fptoui-118.patch"
-      "${patch_dir}/0007-mlir-emitc-Add-support-for-C-API-python-binding-to-E.patch"
-      "${patch_dir}/0008-mlir-emitc-DCE-unimplemented-decls-121253.patch"
-      "${patch_dir}/0009-Re-introduce-Type-Conversion-on-EmitC-121476.patch"
-      "${patch_dir}/0010-mlir-emitc-Fix-invalid-syntax-in-example-of-emitc.re.patch"
-      "${patch_dir}/0011-mlir-emitc-Don-t-emit-extra-semicolon-after-bracket-.patch"
-      "${patch_dir}/0012-mlir-emitc-Expose-emitc-dialect-types-119645.patch"
-      "${patch_dir}/0013-mlir-emitc-Support-convert-arith.extf-and-arith.trun.patch"
-      "${patch_dir}/0014-EmitC-Allow-arrays-of-size-zero-123292.patch"
-      "${patch_dir}/0015-mlir-EmitC-Add-MathToEmitC-pass-for-math-function-lo.patch"
-      "${patch_dir}/0016-mlir-emitc-Set-default-dialect-to-emitc-in-ops-with-.patch"
-      "${patch_dir}/0017-mlir-emitc-Fix-two-EmitC-bugs.patch"
+      "${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}/0008-MLIR-Remove-unnecessary-include-from-MathToEmitC.h-t.patch"
+      "${mlir_patch_dir}/0009-mlir-Support-FileLineColRange-in-LLVM-debug-translat.patch"
+      "${mlir_patch_dir}/0010-MLIR-Fix-LLVMIRTransforms-build-failure-125485.patch"
     # Set the CPM cache key to the Git hash for easy navigation.
     PRE_ADD_HOOK [[
       list(APPEND _vap_UNPARSED_ARGUMENTS
@@ -109,14 +98,18 @@ nv_register_package(
 #-------------------------------------------------------------------------------------
 # Stablehlo
 #-------------------------------------------------------------------------------------
+set(stablehlo_patch_dir "${CMAKE_SOURCE_DIR}/build_tools/patches/stablehlo")
+
 nv_register_package(
   NAME Stablehlo
   VERSION 1.8.0
-  GIT_TAG 6e403b1aa6a71f5eaa09cc720e4ad42f692745e6
+  GIT_TAG 48a1e14edc8219577fcad53de1924876f855f431
   GIT_REPOSITORY "https://github.com/openxla/stablehlo.git"
   PATCHES
-    "${CMAKE_CURRENT_LIST_DIR}/build_tools/patches/stablehlo/0001-transforms-Fix-simplification-patterns-for-stablehlo.patch"
-    "${CMAKE_CURRENT_LIST_DIR}/build_tools/patches/stablehlo/0002-Fix-a-couple-missing-checks-for-static-shapes-in-sta.patch"
+    "${stablehlo_patch_dir}/0001-Fix-a-couple-missing-checks-for-static-shapes-in-sta.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}/0002-cmake-Update-usage-of-HandleLLVMOptions-and-LLVM_DEF.patch"
   OPTIONS
     "STABLEHLO_ENABLE_BINDINGS_PYTHON ON"
     "STABLEHLO_BUILD_EMBEDDED ON"
@@ -163,11 +156,13 @@ nv_register_package(
 #-------------------------------------------------------------------------------------
 # Torch-MLIR
 #-------------------------------------------------------------------------------------
+set(torch_mlir_patch_dir "${CMAKE_SOURCE_DIR}/build_tools/patches/torch_mlir")
+
 nv_register_package(
   NAME torch_mlir
   GIT_REPOSITORY https://github.com/llvm/torch-mlir.git
-  GIT_TAG 30c519369ed7eabad0282d0f874500a9b41fcbbd
-  PATCHES "${CMAKE_CURRENT_LIST_DIR}/build_tools/patches/torch_mlir/torch_mlir.patch"
+  GIT_TAG 169032010793ee7fe3e305ab920e4119fdfc3b11
+  PATCHES "${torch_mlir_patch_dir}/0001-cmake-Allow-finding-Stablehlo-via-find_package.patch"
   EXCLUDE_FROM_ALL TRUE
   # We need to specify an existing directory that is not actually a submodule
   # since GIT_SUBMODULES does not except the empty string due to

mlir-tensorrt/build_tools/patches/mlir/0003-mlir-EmitC-memref-to-emitc-insert-conversion_casts-1.patch

@@ -0,0 +1,41 @@
+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
+
+---
+ 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
+--- 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> {
+         }
+       } // else dim.getIndex is a block argument to reshape->getBlock and
+         // dominates reshape
+-    }   // Check condition 2
++    } // Check condition 2
+     else if (dim->getBlock() != reshape->getBlock() &&
+              !dim.getIndex().getParentRegion()->isProperAncestor(
+                  reshape->getParentRegion())) {
+@@ -1607,9 +1607,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)) {
+-      Type initType = elementsAttr.getType();
+-      Type tensorType = getTensorTypeFromMemRefType(memrefType);
+-      if (initType != tensorType)
++      ShapedType initType = cast<ShapedType>(elementsAttr.getType());
++      ShapedType tensorType =
++          cast<ShapedType>(getTensorTypeFromMemRefType(memrefType));
++      if (initType.getShape() != tensorType.getShape() ||
++          initType.getElementType() != tensorType.getElementType())
+         return emitOpError("initial value expected to be of type ")
+                << tensorType << ", but was of type " << initType;
+     }
+-- 
+2.46.0
+