Skip to content

Commit 0c9f6a0

Browse files
Integrate internal changes.
This PR integrates the following internal changes into OSS. commit c6f9d4b57ae793b08ddf415529a83e2aca6df5ce Author: Sagar Shelke <[email protected]> Date: Tue Jun 10 14:27:27 2025 -0700 [Dialect/Plan] Add constant foldable subgraph analysis pass. This MR adds a generic pass for constant foldable subgraph analysis pass. This pass implements the following three things, 1. A forward data flow analysis to find constant foldable ops. 2. Use clustering to find constant foldable subgraphs. 3. Outline constant foldable subgraphs to `func.func` with `plan.constant_foldable` attributes. Pass option `skipClustering` allows user to extend default pass behavior and skip certain ops from outlining. MLIR test cases are added including region ops tests. commit 5fad6480597ee15bdb4a8c04c033b7ef6beff2d0 Author: Christopher Bate <[email protected]> Date: Tue Jun 10 19:27:04 2025 +0000 NFC: remove unused functions commit 7eb853f52525756e399a7fc6fdc4ab25244d6f94 Author: Chris Bate <[email protected]> Date: Tue Jun 10 12:06:51 2025 -0700 [compiler] Create InferTensorValueRangeInterface Previously, our TensorValueBoundsAnalysis was only able to analyze values that were either constant or were produced by the result of `plan.with_values`. This worked because we rely on `plan.with_values` to link integer range analysis and the TensorValueBoundsAnalysis. However, after clustering there can occur edge cases where the `plan.with_values` op is not present. Furthermore, eventually we would like to avoid creating `plan.with_values` ops in the first place since they temporarily bloat the IR. This commit adds a new interface, InferTensorValueRangeInterface, that allows operations to participate in the TensorValueBoundsAnalysis. The logic for `plan.with_values` is implemented using the interface, and this change adds an implementation and test for `stablehlo.convert`. Besides this new behavior for `stablehlo.convert`, this change is a pure refactor and no other new behavior is added. commit d49fe69e08d79db07bf9f801427c97f6bfbdfcda Author: Christopher Bate <[email protected]> Date: Fri Jun 6 17:03:45 2025 +0000 NFC: Move `python` under `integrations/python` Moves the `python` directory under `integrations/python` to better reflect that, like PJRT, the Python packages are a downstream consumer of the MLIR-TensorRT compiler/runtime libraries. commit ef5490368b2235e80bb21a6aca39db43d24f1c47 Author: Sagar Shelke <[email protected]> Date: Sat Jun 7 14:27:21 2025 -0700 [tensorrt] Implement inliner inferface for TensorRT dialect This MR implements `DialectInlinerInterface` for the TensorRT dialect. All ops in TensorRT dialect are pure and thus can be inlined, except `module`, `call` and `call_alloc`. Upstream `--inline` pass doesn't inline operation with callop interface (`call` and `call_alloc`) because these operations in TensorRT dialect are referring to callee in different symbol table comapred to their own. MLIR valid and invalid test is added. commit cd56aa6a511e2091fcd86106f20d27ff3673db75 Author: Christopher Bate <[email protected]> Date: Wed Jun 11 14:26:03 2025 +000 Fix build with BUILD_SHARED_LIBS=ON The new InferTensorValueRangeInterface was used without correctly specifying the library dependency the PlanIR and StablehloExtIR libraries. GitOrigin-RevId: cd56aa6a511e2091fcd86106f20d27ff3673db75
1 parent 30f9816 commit 0c9f6a0

File tree

62 files changed

+1836
-356
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

62 files changed

+1836
-356
lines changed

mlir-tensorrt/CMakeLists.txt

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -249,4 +249,4 @@ include_directories(${CMAKE_CURRENT_LIST_DIR}/tensorrt/include)
249249
include_directories(${CMAKE_CURRENT_BINARY_DIR}/tensorrt/include)
250250

251251
add_subdirectory(compiler)
252-
add_subdirectory(python)
252+
add_subdirectory(integrations)

mlir-tensorrt/build_tools/docker/Dockerfile

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -74,8 +74,8 @@ EOF
7474
ARG PYTHON_VERSION=3.10
7575
ENV PYENV_ROOT="/pyenv"
7676
ENV PATH="/pyenv/bin:/pyenv/shims:$PATH"
77-
COPY python/requirements-dev.txt /tmp/requirements-dev.txt
78-
COPY python/requirements.txt /tmp/requirements.txt
77+
COPY integrations/python/requirements-dev.txt /tmp/requirements-dev.txt
78+
COPY integrations/python/requirements.txt /tmp/requirements.txt
7979
RUN <<EOF
8080
set -e
8181
case "${LINUX_DISTRO}" in

mlir-tensorrt/build_tools/scripts/build_wheels.sh

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -9,7 +9,7 @@ py_version=${PY_VERSION:-3.10}
99
PATH=$PATH:/pyenv/bin
1010
mkdir -p .private.wheels || true
1111
pyenv local ${py_version}
12-
python${py_version} -m pip install -r python/requirements-dev.txt
12+
python${py_version} -m pip install -r integrations/python/requirements-dev.txt
1313

1414
export DOWNLOAD_TENSORRT_VERSION=${DOWNLOAD_TENSORRT_VERSION:-10.9}
1515

mlir-tensorrt/build_tools/scripts/cicd_build.sh

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -10,7 +10,7 @@ export ENABLE_ASAN=${ENABLE_ASAN:-OFF}
1010
export CPM_SOURCE_CACHE=${CPM_SOURCE_CACHE:-/.cache.cpm}
1111
export CCACHE_DIR=${CCACHE_DIR:-/ccache}
1212

13-
python3 -m pip install -r python/requirements-dev.txt
13+
python3 -m pip install -r integrations/python/requirements-dev.txt
1414

1515
ccache --zero-stats || true
1616
rm -rf ${BUILD_DIR} || true

mlir-tensorrt/compiler/include/mlir-tensorrt/Dialect/Plan/Analysis/BoundsAnalysis.h

Lines changed: 2 additions & 85 deletions
Original file line numberDiff line numberDiff line change
@@ -25,96 +25,13 @@
2525
#ifndef MLIR_TENSORRT_DIALECT_PLAN_ANALYSIS_BOUNDSANALYSIS
2626
#define MLIR_TENSORRT_DIALECT_PLAN_ANALYSIS_BOUNDSANALYSIS
2727

28+
#include "mlir-tensorrt/Interfaces/InferTensorValueRangeInterface.h"
2829
#include "mlir/Analysis/DataFlow/IntegerRangeAnalysis.h"
2930
#include "mlir/Analysis/DataFlow/SparseAnalysis.h"
30-
#include "mlir/Interfaces/InferIntRangeInterface.h"
31-
#include "llvm/Support/raw_ostream.h"
3231

3332
namespace mlir::plan {
3433

35-
//===----------------------------------------------------------------------===//
36-
// BoundsArray
37-
//===----------------------------------------------------------------------===//
38-
39-
/// A BoundsArray is simply an array of ConstantIntRanges used to represent
40-
/// either the bounds on a shape of a tensor-typed SSA value or the bounds
41-
/// of the element values of a statically shaped integer tensor-typed SSA value.
42-
/// When it is used to represent the bounds for the value of a tensor, we use
43-
/// a canonical packed generalized row-major layout mapping from tensor
44-
/// coordinates to storage index.
45-
class BoundsArray {
46-
public:
47-
BoundsArray(
48-
std::optional<SmallVector<ConstantIntRanges>> value = std::nullopt)
49-
: value(std::move(value)) {}
50-
51-
bool isUninitialized() const { return !value.has_value(); }
52-
53-
bool operator==(const BoundsArray &rhs) const { return value == rhs.value; }
54-
55-
ArrayRef<ConstantIntRanges> getValue() const {
56-
assert(!isUninitialized());
57-
return *value;
58-
}
59-
60-
/// Return the most conservative integer scalar bounds for an dynamic/unknown
61-
/// dimension extent.
62-
static ConstantIntRanges getMaxDimRange();
63-
64-
/// Create a BoundsValue from the min/max bounds of shape. Using this method
65-
/// ensures that the `value` are created with the correct storage bitwidth
66-
/// (an implementation detail of the analysis).
67-
static BoundsArray fromShapeBounds(ArrayRef<int64_t> min,
68-
ArrayRef<int64_t> max);
69-
70-
/// Create a `BoundsValue` using the given scalar values encoded as int64_t
71-
/// values. However, when storing the bounds, use the given bitwidth.
72-
/// TODO: remove this when we migrate away from using
73-
/// `#tensorrt.shape_profile` for value bounds.
74-
static BoundsArray fromIntegerValueBounds(unsigned bitwidth,
75-
ArrayRef<int64_t> min,
76-
ArrayRef<int64_t> max);
77-
static BoundsArray fromIntegerValueBounds(ArrayRef<llvm::APInt> min,
78-
ArrayRef<llvm::APInt> max);
79-
80-
/// For the given tensor-typed value, return the most conservative bounds for
81-
/// the shape of `v`. For each unknown dimension of the shape of `v` the
82-
/// `getMaxDimRange()` bound is used.
83-
static BoundsArray getMaxRangeForShapeBounds(Value v);
84-
85-
/// For the given statically shaped integer tensor-typed value, return the
86-
/// most conservative bounds for the value of `v`.
87-
static BoundsArray getMaxRangeForValueBounds(Value v);
88-
89-
/// For the given DenseIntElementsAttr, return a corresponding BoudnsValue
90-
/// representing constant bounds as indicated by the attribute.
91-
static BoundsArray getFromConstantValue(DenseIntElementsAttr attr);
92-
93-
/// Join two BoundsValues by performing a pointwise union of the integer
94-
/// scalar a ranges.
95-
static BoundsArray join(const BoundsArray &lhs, const BoundsArray &rhs);
96-
97-
/// Meet two BoundsValues by performing a pointwise intersection of the
98-
/// integer scalar a ranges.
99-
static BoundsArray meet(const BoundsArray &lhs, const BoundsArray &rhs);
100-
101-
/// Print a human-readable representation of the bounds.
102-
void print(raw_ostream &os) const;
103-
104-
/// Return the min/max bounds representation as two DenseElementsAttrs.
105-
std::pair<DenseElementsAttr, DenseElementsAttr>
106-
getAsElementsAttr(RankedTensorType type) const;
107-
108-
/// Returns DenseElementsAttr representation if the element ranges are all
109-
/// constant (single-value) ranges, otherwise nullopt.
110-
std::optional<DenseElementsAttr>
111-
getConstantValues(RankedTensorType type) const;
112-
113-
private:
114-
std::optional<SmallVector<ConstantIntRanges>> value;
115-
};
116-
117-
llvm::raw_ostream &operator<<(llvm::raw_ostream &os, const BoundsArray &v);
34+
using BoundsArray = mlirtrt::compiler::BoundsArray;
11835

11936
//===----------------------------------------------------------------------===//
12037
// Shape Bounds Analyses

mlir-tensorrt/compiler/include/mlir-tensorrt/Dialect/Plan/IR/CMakeLists.txt

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -17,4 +17,6 @@ add_public_tablegen_target(MLIRTensorRTPlanDialectAttributesIncGen)
1717
set(LLVM_TARGET_DEFINITIONS PlanInterfaces.td)
1818
mlir_tablegen(PlanAttrInterfaces.h.inc -gen-attr-interface-decls)
1919
mlir_tablegen(PlanAttrInterfaces.cpp.inc -gen-attr-interface-defs)
20+
mlir_tablegen(PlanOpInterfaces.h.inc -gen-op-interface-decls)
21+
mlir_tablegen(PlanOpInterfaces.cpp.inc -gen-op-interface-defs)
2022
add_public_tablegen_target(MLIRTensorRTPlanDialectAttrInterfacesIncGen)

mlir-tensorrt/compiler/include/mlir-tensorrt/Dialect/Plan/IR/Plan.h

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -27,6 +27,7 @@
2727
#include "mlir-tensorrt-dialect/Interface/TensorKindOpInterface.h"
2828
#include "mlir-tensorrt/Compiler/Extension.h"
2929
#include "mlir-tensorrt/Dialect/Plan/IR/PlanInterfaces.h"
30+
#include "mlir-tensorrt/Interfaces/InferTensorValueRangeInterface.h"
3031
#include "mlir/Bytecode/BytecodeOpInterface.h"
3132
#include "mlir/IR/BuiltinAttributes.h"
3233
#include "mlir/IR/BuiltinOps.h"

mlir-tensorrt/compiler/include/mlir-tensorrt/Dialect/Plan/IR/PlanOps.td

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -8,6 +8,7 @@ include "mlir/Interfaces/SideEffectInterfaces.td"
88
include "mlir/Interfaces/ControlFlowInterfaces.td"
99
include "mlir/Interfaces/DestinationStyleOpInterface.td"
1010
include "mlir/Interfaces/InferTypeOpInterface.td"
11+
include "mlir-tensorrt/Interfaces/InferTensorValueRangeInterface.td"
1112
include "mlir/IR/OpAsmInterface.td"
1213

1314
class Plan_NativeOpTrait<string name,
@@ -433,6 +434,8 @@ def Plan_WithShapeOp : Plan_Op<"with_shape",
433434
def Plan_WithValuesOp : Plan_Op<"with_values",
434435
[Pure,
435436
DeclareOpInterfaceMethods<TensorKindOpInterface>,
437+
DeclareOpInterfaceMethods<InferTensorValueRangeInterface,
438+
["inferResultRangesFromOptional"]>,
436439
AllTypesMatch<["operand", "result"]>]> {
437440
let summary =
438441
"Ties a tensor value with index SSA values representing its element values";

mlir-tensorrt/compiler/include/mlir-tensorrt/Dialect/Plan/Transforms/Passes.td

Lines changed: 32 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -258,9 +258,9 @@ def ClusteringPass : Pass<"plan-clustering", "::mlir::ModuleOp"> {
258258
operations will be compiled.
259259

260260
The kinds of clusters that can be formed and the specific rules for
261-
clustering are defined by the clustering configuration specified
261+
clustering are defined by the clustering configuration specified
262262
by the module's `plan.cluster_kinds` attribute. This is an array of
263-
attributes which all implement the
263+
attributes which all implement the
264264
[ClusterKindAttrInterface](../IR/PlanInterfaces.td).
265265
}];
266266

@@ -585,5 +585,35 @@ def PlanOwnershipBasedBufferDeallocationPass : Pass<
585585
];
586586
}
587587

588+
//===----------------------------------------------------------------------===//
589+
// PlanOutlineConstantFoldableSubgraphs
590+
//===----------------------------------------------------------------------===//
591+
592+
def PlanOutlineConstantFoldableSubgraphsPass : Pass<
593+
"plan-outline-constant-foldable-subgraphs",
594+
"::mlir::ModuleOp"> {
595+
let summary = "Analyze and outline constant foldable subgraphs";
596+
597+
let description = [{
598+
This pass implements forward dataflow analysis (named `SparseConstantFoldabilityAnalysis`)
599+
to find out constant foldable ops. This analysis, unlike upstream
600+
`ConstantPropagationAnalysis` is very simple and works only for pure ops.
601+
If all operands of an operation are constant foldable, all results are marked
602+
as constant foldable.
603+
Constant foldability analysis is then used along with clustering to
604+
find constant foldable subgraphs. These constant foldable subgraphs are
605+
finally outlined to a private function with `plan.constant_foldable` attribute.
606+
}];
607+
608+
let options = [
609+
Option<"skipClustering", "skip-clustering",
610+
"std::function<bool(Operation*)>", /*default=*/"nullptr",
611+
"This option enables user to extend default pass behavior and skip "
612+
"more ops from clustering. If this method returns true, `op` is not "
613+
"clustered. When op is not clustered, it is not outlined for constant "
614+
"folding. This is helpful in avoiding clustering of ops that can't be "
615+
"run e2e at compile time, in the workflow of user's choice.">,
616+
];
617+
}
588618

589619
#endif // MLIR_TENSORRT_DIALECT_PLAN_TRANSFORMS_PASSES_TD

mlir-tensorrt/compiler/include/mlir-tensorrt/Dialect/StablehloExt/IR/StableHloExt.h

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -34,6 +34,10 @@ void registerTensorKindOpInterfaceExternalModels(DialectRegistry &registry);
3434
/// Register StableHlo op implementations for ReifyRankedShapedTypeOpInterface.
3535
void registerTypeInferenceExternalModels(DialectRegistry &registry);
3636

37+
/// Register StableHlo op implementations for InferTensorValueRangeInterface.
38+
void registerInferTensorValueRangeInterfaceExternalModels(
39+
DialectRegistry &registry);
40+
3741
} // namespace mlir::stablehlo
3842

3943
#endif // MLIR_TENSORRT_DIALECT_STABLEHLOEXT_IR_STABLEHLOEXT_H

0 commit comments

Comments
 (0)