Adding compilation reentrancy tests and new HAL pipeline phases. (#12503)
Two new `--compile-to=` phases are supported:
- `executable-sources`: run just past interface materialization where
`hal.executable` ops with target configurations are present.
- `executable-targets`: run just past executable translation where
`hal.executable.variant` ops have been lowered to their final MLIR form
before linking and serialization (LLVM dialect, SPIR-V dialect, etc).
New tests are added that demonstrate and verify that iree-opt pipelines
can be run piecewise to produce a final output and that `--compile-to=`
at any particular phase can be passed back in and lowered down to a
final output.
A few passes were tweaked to ensure these tests pass.
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertCommandBufferOps.cpp b/compiler/src/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertCommandBufferOps.cpp
index f66bbfc..b68e096 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertCommandBufferOps.cpp
+++ b/compiler/src/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertCommandBufferOps.cpp
@@ -256,11 +256,13 @@
}
return null;
};
+ auto i32Type = rewriter.getI32Type();
+ auto i64Type = rewriter.getI64Type();
SmallVector<Value, 8> callOperands = {
adaptor.getCommandBuffer(),
adaptor.getPipelineLayout(),
- adaptor.getSet(),
+ castToImportType(adaptor.getSet(), i32Type, rewriter),
};
SmallVector<int16_t, 5> segmentSizes = {
/*command_buffer=*/-1,
@@ -270,7 +272,8 @@
static_cast<int16_t>(adaptor.getBindingOrdinals().size()),
};
for (size_t i = 0; i < adaptor.getBindingOrdinals().size(); ++i) {
- callOperands.push_back(adaptor.getBindingOrdinals()[i]);
+ callOperands.push_back(
+ castToImportType(adaptor.getBindingOrdinals()[i], i32Type, rewriter));
auto bindingBuffer = adaptor.getBindingBuffers()[i];
if (bindingBuffer.getType().isa<IREE::VM::RefType>()) {
// Buffer binding; pass 0 for table slot.
@@ -281,10 +284,10 @@
callOperands.push_back(bindingBuffer);
callOperands.push_back(getNull());
}
- callOperands.push_back(castToImportType(adaptor.getBindingOffsets()[i],
- rewriter.getI64Type(), rewriter));
- callOperands.push_back(castToImportType(adaptor.getBindingLengths()[i],
- rewriter.getI64Type(), rewriter));
+ callOperands.push_back(
+ castToImportType(adaptor.getBindingOffsets()[i], i64Type, rewriter));
+ callOperands.push_back(
+ castToImportType(adaptor.getBindingLengths()[i], i64Type, rewriter));
}
auto callOp = rewriter.replaceOpWithNewOp<IREE::VM::CallVariadicOp>(
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Transforms/MaterializeResourceCaches.cpp b/compiler/src/iree/compiler/Dialect/HAL/Transforms/MaterializeResourceCaches.cpp
index 5e485e9..6633df9 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Transforms/MaterializeResourceCaches.cpp
+++ b/compiler/src/iree/compiler/Dialect/HAL/Transforms/MaterializeResourceCaches.cpp
@@ -53,7 +53,33 @@
if (moduleOp.getBody()->empty()) return;
moduleBuilder = OpBuilder(&moduleOp.getBody()->front());
+ // Find all relevant ops. If we don't find any we skip the pass as it's
+ // likely it's already been run. We could fix the pass to better support
+ // partial materialization but there's no use cases for that today.
auto executableOps = llvm::to_vector<8>(moduleOp.getOps<ExecutableOp>());
+ SmallVector<IREE::HAL::DescriptorSetLayoutLookupOp>
+ descriptorSetLayoutLookupOps;
+ SmallVector<IREE::HAL::PipelineLayoutLookupOp> pipelineLayoutLookupOps;
+ SmallVector<IREE::HAL::ExecutableLookupOp> executableLookupOps;
+ for (Operation &funcLikeOp : moduleOp.getOps()) {
+ auto funcOp = llvm::dyn_cast<FunctionOpInterface>(funcLikeOp);
+ if (!funcOp) continue;
+ for (auto &block : funcOp.getFunctionBody()) {
+ block.walk([&](Operation *op) {
+ if (auto lookupOp = dyn_cast<DescriptorSetLayoutLookupOp>(op)) {
+ descriptorSetLayoutLookupOps.push_back(lookupOp);
+ } else if (auto lookupOp = dyn_cast<PipelineLayoutLookupOp>(op)) {
+ pipelineLayoutLookupOps.push_back(lookupOp);
+ } else if (auto lookupOp = dyn_cast<ExecutableLookupOp>(op)) {
+ executableLookupOps.push_back(lookupOp);
+ }
+ });
+ }
+ }
+ if (descriptorSetLayoutLookupOps.empty() &&
+ pipelineLayoutLookupOps.empty() && executableLookupOps.empty()) {
+ return;
+ }
// Declare all layouts used by the executables. This will ensure that the
// initialization order is correct as any pipeline layout needed (and its
@@ -73,28 +99,19 @@
// Declare executable variables so that we can reference them during lookup
// replacement.
for (auto executableOp : executableOps) {
- if (!defineExecutableOp(executableOp)) {
- signalPassFailure();
- return;
- }
+ defineExecutableOp(executableOp);
}
// Generate cached resource singletons and replace lookup ops with direct
// loads from variables.
- for (Operation &funcLikeOp : moduleOp.getOps()) {
- auto funcOp = llvm::dyn_cast<FunctionOpInterface>(funcLikeOp);
- if (!funcOp) continue;
- for (auto &block : funcOp.getFunctionBody()) {
- block.walk([&](Operation *op) {
- if (auto lookupOp = dyn_cast<DescriptorSetLayoutLookupOp>(op)) {
- replaceDescriptorSetLayoutLookupOp(lookupOp);
- } else if (auto lookupOp = dyn_cast<PipelineLayoutLookupOp>(op)) {
- replacePipelineLayoutLookupOp(lookupOp);
- } else if (auto lookupOp = dyn_cast<ExecutableLookupOp>(op)) {
- replaceExecutableLookupOp(lookupOp);
- }
- });
- }
+ for (auto lookupOp : descriptorSetLayoutLookupOps) {
+ replaceDescriptorSetLayoutLookupOp(lookupOp);
+ }
+ for (auto lookupOp : pipelineLayoutLookupOps) {
+ replacePipelineLayoutLookupOp(lookupOp);
+ }
+ for (auto lookupOp : executableLookupOps) {
+ replaceExecutableLookupOp(lookupOp);
}
}
@@ -182,7 +199,7 @@
return globalOp;
}
- IREE::Util::GlobalOp defineExecutableOp(ExecutableOp executableOp) {
+ void defineExecutableOp(ExecutableOp executableOp) {
auto loc = executableOp.getLoc();
auto symbolName =
(StringRef("_executable_") + executableOp.getSymName()).str();
@@ -259,8 +276,6 @@
blockBuilder.create<IREE::Util::GlobalStoreOp>(loc, executableValue,
globalOp.getName());
blockBuilder.create<IREE::Util::InitializerReturnOp>(loc);
-
- return globalOp;
}
// Inlines a constant block as a function in |moduleBuilder| and then inserts
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Transforms/MemoizeDeviceQueries.cpp b/compiler/src/iree/compiler/Dialect/HAL/Transforms/MemoizeDeviceQueries.cpp
index 3f7daf8..5a474ff 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Transforms/MemoizeDeviceQueries.cpp
+++ b/compiler/src/iree/compiler/Dialect/HAL/Transforms/MemoizeDeviceQueries.cpp
@@ -64,6 +64,7 @@
}
// Create each query variable and replace the uses with loads.
+ SymbolTable symbolTable(moduleOp);
auto moduleBuilder = OpBuilder::atBlockBegin(moduleOp.getBody());
for (auto queryKey : llvm::enumerate(deviceQueryKeys)) {
auto queryOps = deviceQueryOps[queryKey.value()];
@@ -82,10 +83,12 @@
auto valueGlobalOp = moduleBuilder.create<IREE::Util::GlobalOp>(
fusedLoc, variableName,
/*isMutable=*/false, queryType);
+ symbolTable.insert(valueGlobalOp);
valueGlobalOp.setPrivate();
auto okGlobalOp = moduleBuilder.create<IREE::Util::GlobalOp>(
fusedLoc, variableName + "_ok",
/*isMutable=*/false, moduleBuilder.getI1Type());
+ symbolTable.insert(okGlobalOp);
okGlobalOp.setPrivate();
auto initializerOp =
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Transforms/Passes.cpp b/compiler/src/iree/compiler/Dialect/HAL/Transforms/Passes.cpp
index 80c20d0..3b555f0 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Transforms/Passes.cpp
+++ b/compiler/src/iree/compiler/Dialect/HAL/Transforms/Passes.cpp
@@ -207,7 +207,8 @@
void buildHALTransformPassPipeline(OpPassManager &passManager,
const TargetOptions &targetOptions,
- const TransformOptions &transformOptions) {
+ const TransformOptions &transformOptions,
+ PipelinePhase compileTo) {
//----------------------------------------------------------------------------
// Device assignment and interface materialization
//----------------------------------------------------------------------------
@@ -225,6 +226,8 @@
createPreprocessExecutablesPass(command));
}
+ if (compileTo == PipelinePhase::ExecutableSources) return;
+
// TODO(benvanik): move translation after conversion; today translation
// inserts the workgroup count logic we need to convert but we could instead
// insert placeholder ops that are expanded after translation.
@@ -238,6 +241,8 @@
passManager.addNestedPass<IREE::HAL::ExecutableOp>(
createTranslateExecutablesPass());
+ if (compileTo == PipelinePhase::ExecutableTargets) return;
+
// Substitute hal.executables we've translated with those specified on the
// command line. This developer feature allows for splicing in hand-authored
// or hand-modified executables in various forms without modifying the
@@ -370,9 +375,11 @@
}
void buildHALTransformPassPipeline(OpPassManager &passManager,
- const TargetOptions &targetOptions) {
+ const TargetOptions &targetOptions,
+ PipelinePhase compileTo) {
TransformOptions transformOptions;
- buildHALTransformPassPipeline(passManager, targetOptions, transformOptions);
+ buildHALTransformPassPipeline(passManager, targetOptions, transformOptions,
+ compileTo);
}
void registerHALConfigurationPassPipeline() {
@@ -390,8 +397,9 @@
"iree-hal-transformation-pipeline",
"Runs the full IREE HAL dialect transformation pipeline",
[](OpPassManager &passManager, const TransformOptions &transformOptions) {
- buildHALTransformPassPipeline(
- passManager, TargetOptions::FromFlags::get(), transformOptions);
+ buildHALTransformPassPipeline(passManager,
+ TargetOptions::FromFlags::get(),
+ transformOptions, PipelinePhase::End);
});
}
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Transforms/Passes.h b/compiler/src/iree/compiler/Dialect/HAL/Transforms/Passes.h
index bcd949e..3c6cdf2 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Transforms/Passes.h
+++ b/compiler/src/iree/compiler/Dialect/HAL/Transforms/Passes.h
@@ -25,6 +25,15 @@
// Helpers
//===----------------------------------------------------------------------===//
+enum class PipelinePhase {
+ // Runs the transform pipeline up to executable sources (pre translation).
+ ExecutableSources,
+ // Runs the transform pipeline until just after executable translation.
+ ExecutableTargets,
+ // Runs the full pipeline.
+ End,
+};
+
// Adds a set of passes to the given pass manager that run the required HAL
// transforms in the canonical order.
//
@@ -35,8 +44,9 @@
// <run conversion to flow/sequencer/etc>
// buildHALTransformPassPipeline & run
// <run conversion from HAL to vm/etc>
-void buildHALTransformPassPipeline(OpPassManager &passManager,
- const TargetOptions &targetOptions);
+void buildHALTransformPassPipeline(
+ OpPassManager &passManager, const TargetOptions &targetOptions,
+ PipelinePhase compileTo = PipelinePhase::End);
// Adds a set of passes to the given pass manager that run the head of the HAL
// pipeline to assign devices, materialize interfaces, and translate
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Transforms/test/materialize_resource_caches.mlir b/compiler/src/iree/compiler/Dialect/HAL/Transforms/test/materialize_resource_caches.mlir
index c11705d..dba9d91 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Transforms/test/materialize_resource_caches.mlir
+++ b/compiler/src/iree/compiler/Dialect/HAL/Transforms/test/materialize_resource_caches.mlir
@@ -208,3 +208,73 @@
}
}
+
+// -----
+
+// Tests that materialization no-ops when resource caches have already been
+// materialized. Today this is rather simplistic and just bails if the names
+// match with the expectation being that users are mostly just running through
+// with --compile-to=hal and not trying to mutate intermediate HAL state. We
+// could rework the pass to support only materializing what's required based on
+// what resources are looked up.
+
+#pipeline_layout_0 = #hal.pipeline.layout<push_constants = 0, sets = [
+ #hal.descriptor_set.layout<0, bindings = [
+ #hal.descriptor_set.binding<0, storage_buffer>,
+ #hal.descriptor_set.binding<1, storage_buffer>
+ ]>
+]>
+
+module attributes {hal.device.targets = [#hal.device.target<"llvm-cpu">]} {
+
+util.global private @_descriptor_set_layout_0 : !hal.descriptor_set_layout
+util.initializer {
+ %device = hal.ex.shared_device : !hal.device
+ %descriptor_set_layout = hal.descriptor_set_layout.create device(%device : !hal.device) flags("None") bindings([#hal.descriptor_set.binding<0, storage_buffer>, #hal.descriptor_set.binding<1, storage_buffer>]) : !hal.descriptor_set_layout
+ util.global.store %descriptor_set_layout, @_descriptor_set_layout_0 : !hal.descriptor_set_layout
+ util.initializer.return
+}
+
+util.global private @_pipeline_layout_0 : !hal.pipeline_layout
+util.initializer {
+ %_descriptor_set_layout_0 = util.global.load @_descriptor_set_layout_0 : !hal.descriptor_set_layout
+ %device = hal.ex.shared_device : !hal.device
+ %pipeline_layout = hal.pipeline_layout.create device(%device : !hal.device) push_constants(0) layouts([%_descriptor_set_layout_0]) : !hal.pipeline_layout
+ util.global.store %pipeline_layout, @_pipeline_layout_0 : !hal.pipeline_layout
+ util.initializer.return
+}
+
+util.global private @_executable_exe : !hal.executable
+util.initializer {
+ %device = hal.ex.shared_device : !hal.device
+ %0 = hal.device.switch<%device : !hal.device> -> !hal.executable
+ #hal.device.match.executable.format<"vmvx-bytecode-fb"> {
+ %_pipeline_layout_0 = util.global.load @_pipeline_layout_0 : !hal.pipeline_layout
+ %exe = hal.executable.create device(%device : !hal.device) target(@exe0::@vmvx) layouts([%_pipeline_layout_0]) : !hal.executable
+ hal.return %exe : !hal.executable
+ },
+ #hal.match.always {
+ %1 = util.null : !hal.executable
+ hal.return %1 : !hal.executable
+ }
+ util.global.store %0, @_executable_exe : !hal.executable
+ util.initializer.return
+}
+
+hal.executable @exe {
+ hal.executable.variant @vmvx, target = <"vmvx", "vmvx-bytecode-fb"> {
+ hal.executable.export @entry ordinal(0) layout(#pipeline_layout_0) attributes {
+ workgroup_size = [32 : index, 1 : index, 1 : index]
+ }
+ }
+}
+
+// CHECK-LABEL: @exeLookup
+func.func @exeLookup(%device : !hal.device) -> !hal.executable {
+ // CHECK: %[[EXE:.+]] = util.global.load @_executable_exe : !hal.executable
+ %0 = util.global.load @_executable_exe : !hal.executable
+ // CHECK-NEXT: return %[[EXE]]
+ return %0 : !hal.executable
+}
+
+}
diff --git a/compiler/src/iree/compiler/Dialect/Stream/Transforms/VerifyLowerings.cpp b/compiler/src/iree/compiler/Dialect/Stream/Transforms/VerifyLowerings.cpp
index 0902d82..b2816b5 100644
--- a/compiler/src/iree/compiler/Dialect/Stream/Transforms/VerifyLowerings.cpp
+++ b/compiler/src/iree/compiler/Dialect/Stream/Transforms/VerifyLowerings.cpp
@@ -303,11 +303,11 @@
void runOnOperation() override {
// We cannot have stream.cmd.* ops mixed with stream.tensor/async.* ops
- // as they use different memory models.
+ // as they use different memory models. We need to allow them through,
+ // though, to allow for compiler re-entrancy.
Verifier verifier;
setupDefaultOpLegality(verifier);
markTensorInputsIllegal(verifier);
- markStreamCmdOpsIllegal(verifier);
if (failed(verifier.run(getOperation()))) {
return signalPassFailure();
}
@@ -339,12 +339,12 @@
void runOnOperation() override {
// We cannot have stream.cmd.* ops mixed with stream.tensor/async.* ops
- // as they use different memory models.
+ // as they use different memory models. We need to allow them through,
+ // though, to allow for compiler re-entrancy.
Verifier verifier;
setupDefaultOpLegality(verifier);
markTensorInputsIllegal(verifier);
markStreamTensorOpsIllegal(verifier);
- markStreamCmdOpsIllegal(verifier);
// All resources should have had their usage assigned.
verifier.addTypeVerifier<IREE::Stream::ResourceType>([](auto type) {
@@ -357,6 +357,11 @@
// All streamable ops should be inside of execution regions.
verifier.addOpVerifier<IREE::Stream::StreamableOpInterface>(
[](auto op) -> Optional<Verifier::Legality> {
+ // Skip cmd ops that may exist.
+ if (op->template hasTrait<OpTrait::IREE::Stream::CmdPhaseOp>()) {
+ return Verifier::Legality::LEGAL;
+ }
+
// Allow metadata ops outside of execution regions.
if (op.isMetadata()) return Verifier::Legality::LEGAL;
diff --git a/compiler/src/iree/compiler/Dialect/Util/IR/UtilOps.td b/compiler/src/iree/compiler/Dialect/Util/IR/UtilOps.td
index 500f69c..9c79984 100644
--- a/compiler/src/iree/compiler/Dialect/Util/IR/UtilOps.td
+++ b/compiler/src/iree/compiler/Dialect/Util/IR/UtilOps.td
@@ -350,7 +350,11 @@
let hasCanonicalizer = 1;
}
-def Util_UnreachableOp : Util_Op<"unreachable", [NoMemoryEffect, Terminator]> {
+def Util_UnreachableOp : Util_Op<"unreachable", [
+ NoMemoryEffect,
+ ReturnLike,
+ Terminator
+ ]> {
let summary = [{unreachable assertion op}];
let description = [{
Signals to the compiler that the parent block should not be reachable.
diff --git a/compiler/src/iree/compiler/Pipelines/Pipelines.cpp b/compiler/src/iree/compiler/Pipelines/Pipelines.cpp
index 904f297..cb1c611 100644
--- a/compiler/src/iree/compiler/Pipelines/Pipelines.cpp
+++ b/compiler/src/iree/compiler/Pipelines/Pipelines.cpp
@@ -145,6 +145,19 @@
break;
}
+ IREE::HAL::PipelinePhase halCompileTo;
+ switch (compileTo) {
+ default:
+ halCompileTo = IREE::HAL::PipelinePhase::End;
+ break;
+ case IREEVMPipelinePhase::ExecutableSources:
+ halCompileTo = IREE::HAL::PipelinePhase::ExecutableSources;
+ break;
+ case IREEVMPipelinePhase::ExecutableTargets:
+ halCompileTo = IREE::HAL::PipelinePhase::ExecutableTargets;
+ break;
+ }
+
IREE_TRACE_ADD_BEGIN_FRAME_PASS(passManager, "HAL");
switch (schedulingOptions.executionModel) {
case SchedulingOptions::ExecutionModel::HostOnly:
@@ -153,7 +166,8 @@
default:
case SchedulingOptions::ExecutionModel::AsyncInternal:
case SchedulingOptions::ExecutionModel::AsyncExternal:
- IREE::HAL::buildHALTransformPassPipeline(passManager, executableOptions);
+ IREE::HAL::buildHALTransformPassPipeline(passManager, executableOptions,
+ halCompileTo);
break;
case SchedulingOptions::ExecutionModel::InlineStatic:
IREE::HAL::Inline::buildHALInlineStaticTransformPassPipeline(
@@ -165,7 +179,10 @@
break;
}
IREE_TRACE_ADD_END_FRAME_PASS(passManager, "HAL");
- if (compileTo == IREEVMPipelinePhase::HAL) return; // early-exit
+ if (compileTo == IREEVMPipelinePhase::HAL ||
+ halCompileTo != IREE::HAL::PipelinePhase::End) {
+ return; // early-exit
+ }
IREE_TRACE_ADD_BEGIN_FRAME_PASS(passManager, "VM");
IREE::VM::buildVMTransformPassPipeline(passManager, targetOptions);
diff --git a/compiler/src/iree/compiler/Pipelines/Pipelines.h b/compiler/src/iree/compiler/Pipelines/Pipelines.h
index c8781b2..c135f47 100644
--- a/compiler/src/iree/compiler/Pipelines/Pipelines.h
+++ b/compiler/src/iree/compiler/Pipelines/Pipelines.h
@@ -34,6 +34,8 @@
Preprocessing,
Flow,
Stream,
+ ExecutableSources,
+ ExecutableTargets,
HAL,
VM,
End,
@@ -54,6 +56,11 @@
"Compiles up to the `flow` dialect.");
callback(IREEVMPipelinePhase::Stream, "stream",
"Compiles up to the `stream` dialect.");
+ callback(IREEVMPipelinePhase::ExecutableSources, "executable-sources",
+ "Compiles up to just before `hal.executable`s are translated, "
+ "excluding codegen.");
+ callback(IREEVMPipelinePhase::ExecutableTargets, "executable-targets",
+ "Compiles up to translated `hal.executable`s, including codegen.");
callback(IREEVMPipelinePhase::HAL, "hal",
"Compiles up to the `hal` dialect, including codegen.");
callback(IREEVMPipelinePhase::VM, "vm", "Compiles up to the `vm` dialect.");
diff --git a/tools/test/BUILD b/tools/test/BUILD
index 8bd9fa6..f7bf4f8 100644
--- a/tools/test/BUILD
+++ b/tools/test/BUILD
@@ -18,6 +18,8 @@
name = "lit",
srcs = enforce_glob(
[
+ "compile_pipelines.mlir",
+ "compile_to_continuation.mlir",
"compile_to_phase.mlir",
"executable_benchmarks.mlir",
"executable_sources.mlir",
@@ -46,6 +48,7 @@
tools = [
"//tools:iree-benchmark-module",
"//tools:iree-compile",
+ "//tools:iree-opt",
"//tools:iree-run-mlir",
"//tools:iree-run-module",
"@llvm-project//lld",
diff --git a/tools/test/CMakeLists.txt b/tools/test/CMakeLists.txt
index 64a3e22..7c92d9d 100644
--- a/tools/test/CMakeLists.txt
+++ b/tools/test/CMakeLists.txt
@@ -14,6 +14,8 @@
NAME
lit
SRCS
+ "compile_pipelines.mlir"
+ "compile_to_continuation.mlir"
"compile_to_phase.mlir"
"executable_benchmarks.mlir"
"executable_sources.mlir"
@@ -32,6 +34,7 @@
FileCheck
iree-benchmark-module
iree-compile
+ iree-opt
iree-run-mlir
iree-run-module
not
diff --git a/tools/test/compile_pipelines.mlir b/tools/test/compile_pipelines.mlir
new file mode 100644
index 0000000..7275ec9
--- /dev/null
+++ b/tools/test/compile_pipelines.mlir
@@ -0,0 +1,14 @@
+// RUN: iree-opt --iree-common-input-transformation-pipeline %s | \
+// RUN: iree-opt --iree-abi-transformation-pipeline - | \
+// RUN: iree-opt --iree-common-input-transformation-pipeline - | \
+// RUN: iree-opt --iree-flow-transformation-pipeline - | \
+// RUN: iree-opt --iree-stream-transformation-pipeline - | \
+// RUN: iree-opt --iree-hal-transformation-pipeline --iree-hal-target-backends=vmvx - | \
+// RUN: iree-opt --iree-vm-transformation-pipeline - | \
+// RUN: FileCheck %s
+
+// CHECK: vm.func private @abs(%arg0: !vm.ref<!hal.buffer_view>) -> !vm.ref<!hal.buffer_view>
+func.func @abs(%input : tensor<f32>) -> (tensor<f32>) {
+ %result = math.absf %input : tensor<f32>
+ return %result : tensor<f32>
+}
diff --git a/tools/test/compile_to_continuation.mlir b/tools/test/compile_to_continuation.mlir
new file mode 100644
index 0000000..631afca
--- /dev/null
+++ b/tools/test/compile_to_continuation.mlir
@@ -0,0 +1,44 @@
+// RUN: iree-compile --compile-to=input %s | \
+// RUN: iree-compile --output-format=vm-asm --iree-hal-target-backends=vmvx - | \
+// RUN: FileCheck %s --check-prefix=INPUT-PHASE
+// INPUT-PHASE: vm.func private @abs(%arg0: !vm.ref<!hal.buffer_view>) -> !vm.ref<!hal.buffer_view>
+
+// RUN: iree-compile --compile-to=abi %s | \
+// RUN: iree-compile --output-format=vm-asm --iree-hal-target-backends=vmvx - | \
+// RUN: FileCheck %s --check-prefix=ABI-PHASE
+// ABI-PHASE: vm.func private @abs(%arg0: !vm.ref<!hal.buffer_view>) -> !vm.ref<!hal.buffer_view>
+
+// RUN: iree-compile --compile-to=flow %s | \
+// RUN: iree-compile --output-format=vm-asm --iree-hal-target-backends=vmvx - | \
+// RUN: FileCheck %s --check-prefix=FLOW-PHASE
+// FLOW-PHASE: vm.func private @abs(%arg0: !vm.ref<!hal.buffer_view>) -> !vm.ref<!hal.buffer_view>
+
+// RUN: iree-compile --compile-to=stream %s | \
+// RUN: iree-compile --output-format=vm-asm --iree-hal-target-backends=vmvx - | \
+// RUN: FileCheck %s --check-prefix=STREAM-PHASE
+// STREAM-PHASE: vm.func private @abs(%arg0: !vm.ref<!hal.buffer_view>) -> !vm.ref<!hal.buffer_view>
+
+// RUN: iree-compile --compile-to=executable-sources --iree-hal-target-backends=vmvx %s | \
+// RUN: iree-compile --output-format=vm-asm - | \
+// RUN: FileCheck %s --check-prefix=EXECUTABLE-SOURCES-PHASE
+// EXECUTABLE-SOURCES-PHASE: vm.func private @abs(%arg0: !vm.ref<!hal.buffer_view>) -> !vm.ref<!hal.buffer_view>
+
+// RUN: iree-compile --compile-to=executable-targets --iree-hal-target-backends=vmvx %s | \
+// RUN: iree-compile --output-format=vm-asm - | \
+// RUN: FileCheck %s --check-prefix=EXECUTABLE-TARGETS-PHASE
+// EXECUTABLE-TARGETS-PHASE: vm.func private @abs(%arg0: !vm.ref<!hal.buffer_view>) -> !vm.ref<!hal.buffer_view>
+
+// RUN: iree-compile --compile-to=hal --iree-hal-target-backends=vmvx %s | \
+// RUN: iree-compile --output-format=vm-asm - | \
+// RUN: FileCheck %s --check-prefix=HAL-PHASE
+// HAL-PHASE: vm.func private @abs(%arg0: !vm.ref<!hal.buffer_view>) -> !vm.ref<!hal.buffer_view>
+
+// RUN: iree-compile --compile-to=vm --iree-hal-target-backends=vmvx %s | \
+// RUN: iree-compile --output-format=vm-asm - | \
+// RUN: FileCheck %s --check-prefix=VM-PHASE
+// VM-PHASE: vm.func private @abs(%arg0: !vm.ref<!hal.buffer_view>) -> !vm.ref<!hal.buffer_view>
+
+func.func @abs(%input : tensor<f32>) -> (tensor<f32>) {
+ %result = math.absf %input : tensor<f32>
+ return %result : tensor<f32>
+}
diff --git a/tools/test/compile_to_phase.mlir b/tools/test/compile_to_phase.mlir
index cb3a8ed..b88dcac 100644
--- a/tools/test/compile_to_phase.mlir
+++ b/tools/test/compile_to_phase.mlir
@@ -1,20 +1,31 @@
-// RUN: iree-compile --compile-to=input --iree-hal-target-backends=vmvx %s | FileCheck %s --check-prefix=INPUT-PHASE
+// RUN: iree-compile --compile-to=input %s | FileCheck %s --check-prefix=INPUT-PHASE
// INPUT-PHASE: func.func @abs(%[[ARG0:.+]]: tensor<f32>)
// INPUT-PHASE: math.absf %[[ARG0]] : tensor<f32>
-// RUN: iree-compile --compile-to=abi --iree-hal-target-backends=vmvx %s | FileCheck %s --check-prefix=ABI-PHASE
+// RUN: iree-compile --compile-to=abi %s | FileCheck %s --check-prefix=ABI-PHASE
// ABI-PHASE: func.func @abs(%[[ARG0:.+]]: !hal.buffer_view)
// ABI-PHASE: %[[INPUT:.+]] = hal.tensor.import %[[ARG0]] : !hal.buffer_view -> tensor<f32>
// ABI-PHASE: math.absf %[[INPUT]] : tensor<f32>
-// RUN: iree-compile --compile-to=flow --iree-hal-target-backends=vmvx %s | FileCheck %s --check-prefix=FLOW-PHASE
+// RUN: iree-compile --compile-to=flow %s | FileCheck %s --check-prefix=FLOW-PHASE
// FLOW-PHASE: flow.executable.export public @abs_dispatch_0
// FLOW-PHASE: flow.dispatch @abs_dispatch_0
-// RUN: iree-compile --compile-to=stream --iree-hal-target-backends=vmvx %s | FileCheck %s --check-prefix=STREAM-PHASE
+// RUN: iree-compile --compile-to=stream %s | FileCheck %s --check-prefix=STREAM-PHASE
// STREAM-PHASE: stream.executable.export public @abs_dispatch_0
// STREAM-PHASE: stream.cmd.dispatch @abs_dispatch_0
+// RUN: iree-compile --compile-to=executable-sources --iree-hal-target-backends=vmvx %s | FileCheck %s --check-prefix=EXECUTABLE-SOURCES-PHASE
+// EXECUTABLE-SOURCES-PHASE: hal.executable private @abs_dispatch_0
+// EXECUTABLE-SOURCES-PHASE: hal.executable.variant
+// EXECUTABLE-SOURCES-PHASE: linalg.generic
+// EXECUTABLE-SOURCES-PHASE: math.absf
+
+// RUN: iree-compile --compile-to=executable-targets --iree-hal-target-backends=vmvx %s | FileCheck %s --check-prefix=EXECUTABLE-TARGETS-PHASE
+// EXECUTABLE-TARGETS-PHASE: hal.executable private @abs_dispatch_0
+// EXECUTABLE-TARGETS-PHASE: hal.executable.variant
+// EXECUTABLE-TARGETS-PHASE: vm.abs.f32
+
// RUN: iree-compile --compile-to=hal --iree-hal-target-backends=vmvx %s | FileCheck %s --check-prefix=HAL-PHASE
// HAL-PHASE: hal.executable private @abs_dispatch_0
// HAL-PHASE: hal.executable.binary
