blob: 2f7cd96ee5f5f31ac8201c6e7932b1571f3c1115 [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
#include "compiler/plugins/input/StableHLO/Conversion/CustomCallAndCompositePreprocessing.h"
#include "compiler/plugins/input/StableHLO/Conversion/StableHLOCustomCallMetadata.h"
#include "compiler/plugins/input/StableHLO/Conversion/TargetRegistry.h"
#include <gtest/gtest.h>
#include <atomic>
#include "mlir/Dialect/Arith/IR/Arith.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/IR/BuiltinOps.h"
#include "mlir/IR/Verifier.h"
#include "mlir/Parser/Parser.h"
#include "mlir/Pass/PassManager.h"
#include "stablehlo/transforms/Passes.h"
namespace mlir::iree_compiler::stablehlo {
namespace {
constexpr StringLiteral kCustomCallModule = R"mlir(
module {
func.func @run() -> tensor<i32> {
%result = stablehlo.custom_call @example.constant() {
api_version = 4 : i32, backend_config = {configuration = "preserved"},
has_side_effect = false
} : () -> tensor<i32>
return %result : tensor<i32>
}
}
)mlir";
class CustomCallAndCompositePreprocessingTest : public ::testing::Test {
protected:
CustomCallAndCompositePreprocessingTest() {
context.loadDialect<func::FuncDialect, mlir::stablehlo::StablehloDialect>();
}
void checkCustomCall(bool serializeToVhlo) {
auto module = parseSourceString<ModuleOp>(kCustomCallModule, &context);
ASSERT_TRUE(module);
if (serializeToVhlo) {
PassManager serialization(&context);
serialization.addPass(
mlir::stablehlo::createStablehloLegalizeToVhloPass());
ASSERT_TRUE(succeeded(serialization.run(*module)));
bool foundVhlo = false;
module->walk([&](Operation *op) {
foundVhlo |= op->getName().getDialectNamespace() == "vhlo";
});
ASSERT_TRUE(foundVhlo);
}
auto calls = std::make_shared<std::atomic<unsigned>>(0);
TargetRegistry registry;
ASSERT_TRUE(succeeded(registry.addCustomCall(
"example.constant",
[calls](mlir::stablehlo::CustomCallOp op,
PatternRewriter &rewriter) -> LogicalResult {
EXPECT_EQ(op->getAttrOfType<DictionaryAttr>("backend_config")
.getAs<StringAttr>("configuration")
.getValue(),
"preserved");
EXPECT_EQ(op->getAttrOfType<IntegerAttr>("api_version").getInt(), 4);
EXPECT_FALSE(hasEmptyCustomCallMetadata(op));
++*calls;
auto type = cast<RankedTensorType>(op.getResult(0).getType());
auto constant = arith::ConstantOp::create(
rewriter, op.getLoc(), type,
DenseIntElementsAttr::get(type, APInt(32, 7)));
rewriter.replaceOp(op, constant.getResult());
return success();
},
[](DialectRegistry &dialects) {
dialects.insert<arith::ArithDialect>();
})));
PassManager preprocessing(&context);
buildCustomCallAndCompositePreprocessingPassPipeline(preprocessing,
registry);
ASSERT_TRUE(succeeded(preprocessing.run(*module)));
EXPECT_EQ(calls->load(), 1u);
EXPECT_TRUE(succeeded(verify(*module)));
module->walk([&](Operation *op) {
EXPECT_NE(op->getName().getDialectNamespace(), "vhlo");
EXPECT_FALSE(isa<mlir::stablehlo::CustomCallOp>(op));
});
auto func = module->lookupSymbol<func::FuncOp>("run");
ASSERT_TRUE(func);
auto result = cast<func::ReturnOp>(func.getBody().front().getTerminator())
.getOperands()
.front();
auto constant = result.getDefiningOp<arith::ConstantOp>();
ASSERT_TRUE(constant);
auto value = cast<DenseIntElementsAttr>(constant.getValue());
EXPECT_EQ(value.getSplatValue<APInt>().getSExtValue(), 7);
}
MLIRContext context;
};
TEST_F(CustomCallAndCompositePreprocessingTest,
HandlesStablehloBeforeInputConversion) {
checkCustomCall(false);
}
TEST_F(CustomCallAndCompositePreprocessingTest,
NormalizesVhloBeforeTargetLookup) {
checkCustomCall(true);
}
TEST_F(CustomCallAndCompositePreprocessingTest,
LeavesOtherInputDialectsUnchanged) {
context.loadDialect<arith::ArithDialect>();
auto module = parseSourceString<ModuleOp>(R"mlir(
module {
func.func @run(%x: i32) -> i32 {
%c = arith.constant 1 : i32
%result = arith.addi %x, %c : i32
return %result : i32
}
}
)mlir",
&context);
ASSERT_TRUE(module);
auto original = module->clone();
TargetRegistry registry;
PassManager preprocessing(&context);
buildCustomCallAndCompositePreprocessingPassPipeline(preprocessing, registry);
ASSERT_TRUE(succeeded(preprocessing.run(*module)));
std::string before, after;
llvm::raw_string_ostream beforeStream(before), afterStream(after);
original->print(beforeStream);
module->print(afterStream);
EXPECT_EQ(before, after);
original->destroy();
}
} // namespace
} // namespace mlir::iree_compiler::stablehlo