blob: 0346bb2a26c875e48e559ccedb9a885572ee1ee8 [file]
// Copyright 2026 The IREE Authors
//
// Licensed 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
// Lowers `#stablehlo.bounds` to a `util.assume.int` on each bounded
// dimension, tied back to the tensor.
#include "compiler/plugins/input/StableHLO/Conversion/Preprocessing/Passes.h"
#include "iree/compiler/Dialect/Flow/IR/FlowDialect.h"
#include "iree/compiler/Dialect/Flow/IR/FlowOps.h"
#include "iree/compiler/Dialect/Util/IR/UtilDialect.h"
#include "iree/compiler/Dialect/Util/IR/UtilOps.h"
#include "llvm/ADT/SmallVectorExtras.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/Dialect/Tensor/IR/Tensor.h"
#include "mlir/IR/AttrTypeSubElements.h"
#include "mlir/IR/BuiltinTypes.h"
#include "stablehlo/dialect/StablehloOps.h"
namespace mlir::iree_compiler::stablehlo {
#define GEN_PASS_DEF_LOWERBOUNDS
#include "compiler/plugins/input/StableHLO/Conversion/Preprocessing/Passes.h.inc"
namespace {
mlir::stablehlo::TypeExtensionsAttr getBounds(Type type) {
auto tensorType = dyn_cast<RankedTensorType>(type);
if (!tensorType) {
return nullptr;
}
return dyn_cast_if_present<mlir::stablehlo::TypeExtensionsAttr>(
tensorType.getEncoding());
}
Type stripBounds(Type type) {
if (!getBounds(type)) {
return type;
}
auto tensorType = cast<RankedTensorType>(type);
return RankedTensorType::get(tensorType.getShape(),
tensorType.getElementType());
}
// Rewrites `value` to its plain type and routes its uses through a
// `flow.tensor.tie_shape` whose dynamic dims carry the bounds.
void applyBounds(OpBuilder &builder, Value value) {
auto bounds = getBounds(value.getType());
if (!bounds) {
return;
}
auto tensorType = cast<RankedTensorType>(value.getType());
auto plainType =
RankedTensorType::get(tensorType.getShape(), tensorType.getElementType());
value.setType(plainType);
if (plainType.hasStaticShape()) {
return;
}
OpBuilder::InsertionGuard guard(builder);
builder.setInsertionPointAfterValue(value);
SmallVector<OpOperand *> originalUses;
for (OpOperand &use : value.getUses()) {
originalUses.push_back(&use);
}
Location loc = value.getLoc();
SmallVector<Value> dynamicDims;
for (auto [index, size] : llvm::enumerate(plainType.getShape())) {
if (!ShapedType::isDynamic(size)) {
continue;
}
auto constOp = arith::ConstantIndexOp::create(builder, loc, index);
auto dimOp = tensor::DimOp::create(builder, loc, value, constOp);
Value dim = dimOp;
int64_t bound = bounds.getBounds()[index];
if (!ShapedType::isDynamic(bound)) {
auto assumption = builder.getAttr<IREE::Util::IntAssumptionAttr>(
/*umin=*/std::nullopt, /*umax=*/static_cast<uint64_t>(bound),
/*udiv=*/std::nullopt);
auto assumeOp =
IREE::Util::AssumeIntOp::create(builder, loc, dim, assumption);
dim = assumeOp.getResult(0);
}
dynamicDims.push_back(dim);
}
auto tieOp = IREE::Flow::TensorTieShapeOp::create(builder, loc, plainType,
value, dynamicDims);
for (OpOperand *use : originalUses) {
use->set(tieOp.getResult());
}
}
struct LowerBounds final : impl::LowerBoundsBase<LowerBounds> {
void runOnOperation() override {
func::FuncOp funcOp = getOperation();
OpBuilder builder(funcOp.getContext());
AttrTypeReplacer constantAttrReplacer;
constantAttrReplacer.addReplacement(
[](DenseElementsAttr value) -> Attribute {
auto plainType = cast<ShapedType>(stripBounds(value.getType()));
if (plainType == value.getType()) {
return value;
}
return value.reshape(plainType);
});
funcOp.walk([&](Block *block) {
for (BlockArgument arg : block->getArguments()) {
applyBounds(builder, arg);
}
});
funcOp.walk([&](Operation *op) {
// Constants infer their result types from their value attributes. Keep
// those types consistent when removing the encoding from the results.
if (op->hasTrait<OpTrait::ConstantLike>()) {
constantAttrReplacer.replaceElementsIn(op, /*replaceAttrs=*/true,
/*replaceLocs=*/false,
/*replaceTypes=*/false);
}
for (Value result : op->getResults()) {
applyBounds(builder, result);
}
});
// The signature is not reached by the walks above.
FunctionType oldType = funcOp.getFunctionType();
funcOp.setType(FunctionType::get(
funcOp.getContext(),
llvm::map_to_vector(oldType.getInputs(), stripBounds),
llvm::map_to_vector(oldType.getResults(), stripBounds)));
}
};
} // namespace
} // namespace mlir::iree_compiler::stablehlo