Adding hal_loader dialect and runtime module.
This pairs with the hal_inline dialect to allow for dynamically loaded
executables just as with the full HAL dialect only with the ability to
execute them inline. By doing this we can use codegen to target native
code for the executables while still running everything fully
synchronously inline on host-local buffers.
diff --git a/compiler/src/iree/compiler/Codegen/Common/TileDispatchUsingInterface.cpp b/compiler/src/iree/compiler/Codegen/Common/TileDispatchUsingInterface.cpp
index 289c3d0..fc28200 100644
--- a/compiler/src/iree/compiler/Codegen/Common/TileDispatchUsingInterface.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/TileDispatchUsingInterface.cpp
@@ -455,7 +455,7 @@
if (failed(tiledProducer)) {
return failure();
}
- rewriter.replaceOp(sliceOp, tiledProducer.getValue());
+ rewriter.replaceOp(sliceOp, tiledProducer.value());
return success();
}
};
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/BUILD b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/BUILD
new file mode 100644
index 0000000..f7ac0a4
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/BUILD
@@ -0,0 +1,22 @@
+# Copyright 2022 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
+
+load("//build_tools/embed_data:build_defs.bzl", "c_embed_data")
+
+package(
+ default_visibility = ["//visibility:public"],
+ features = ["layering_check"],
+ licenses = ["notice"], # Apache 2.0
+)
+
+c_embed_data(
+ name = "hal_loader_imports",
+ srcs = ["hal_loader.imports.mlir"],
+ c_file_output = "hal_loader.imports.c",
+ flatten = True,
+ h_file_output = "hal_loader.imports.h",
+ identifier = "iree_hal_loader_imports",
+)
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/CMakeLists.txt b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/CMakeLists.txt
new file mode 100644
index 0000000..9a3effd
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/CMakeLists.txt
@@ -0,0 +1,28 @@
+################################################################################
+# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from #
+# compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/BUILD #
+# #
+# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary #
+# CMake-only content. #
+# #
+# To disable autogeneration for this file entirely, delete this header. #
+################################################################################
+
+iree_add_all_subdirs()
+
+iree_c_embed_data(
+ NAME
+ hal_loader_imports
+ SRCS
+ "hal_loader.imports.mlir"
+ C_FILE_OUTPUT
+ "hal_loader.imports.c"
+ H_FILE_OUTPUT
+ "hal_loader.imports.h"
+ IDENTIFIER
+ "iree_hal_loader_imports"
+ FLATTEN
+ PUBLIC
+)
+
+### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ###
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/BUILD b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/BUILD
new file mode 100644
index 0000000..cadf491
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/BUILD
@@ -0,0 +1,11 @@
+# Copyright 2022 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
+
+package(
+ default_visibility = ["//visibility:public"],
+ features = ["layering_check"],
+ licenses = ["notice"], # Apache 2.0
+)
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/CMakeLists.txt b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/CMakeLists.txt
new file mode 100644
index 0000000..4bd0760
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/CMakeLists.txt
@@ -0,0 +1,13 @@
+################################################################################
+# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from #
+# compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/BUILD #
+# #
+# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary #
+# CMake-only content. #
+# #
+# To disable autogeneration for this file entirely, delete this header. #
+################################################################################
+
+iree_add_all_subdirs()
+
+### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ###
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/HALLoaderToVM/BUILD b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/HALLoaderToVM/BUILD
new file mode 100644
index 0000000..a89e383
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/HALLoaderToVM/BUILD
@@ -0,0 +1,36 @@
+# Copyright 2022 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
+
+load("//build_tools/bazel:build_defs.oss.bzl", "iree_compiler_cc_library")
+
+package(
+ default_visibility = ["//visibility:public"],
+ features = ["layering_check"],
+ licenses = ["notice"], # Apache 2.0
+)
+
+iree_compiler_cc_library(
+ name = "HALLoaderToVM",
+ srcs = [
+ "ConvertHALLoaderToVM.cpp",
+ ],
+ hdrs = [
+ "ConvertHALLoaderToVM.h",
+ ],
+ deps = [
+ "//compiler/src/iree/compiler/Dialect/HAL/Conversion/HALToVM",
+ "//compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR",
+ "//compiler/src/iree/compiler/Dialect/Util/IR",
+ "//compiler/src/iree/compiler/Dialect/VM/Conversion",
+ "//compiler/src/iree/compiler/Dialect/VM/Conversion/StandardToVM",
+ "//compiler/src/iree/compiler/Dialect/VM/IR",
+ "@llvm-project//mlir:ArithmeticDialect",
+ "@llvm-project//mlir:FuncDialect",
+ "@llvm-project//mlir:IR",
+ "@llvm-project//mlir:Pass",
+ "@llvm-project//mlir:Transforms",
+ ],
+)
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/HALLoaderToVM/CMakeLists.txt b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/HALLoaderToVM/CMakeLists.txt
new file mode 100644
index 0000000..3facc92
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/HALLoaderToVM/CMakeLists.txt
@@ -0,0 +1,35 @@
+################################################################################
+# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from #
+# compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/HALLoaderToVM/BUILD#
+# #
+# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary #
+# CMake-only content. #
+# #
+# To disable autogeneration for this file entirely, delete this header. #
+################################################################################
+
+iree_add_all_subdirs()
+
+iree_cc_library(
+ NAME
+ HALLoaderToVM
+ HDRS
+ "ConvertHALLoaderToVM.h"
+ SRCS
+ "ConvertHALLoaderToVM.cpp"
+ DEPS
+ MLIRArithmeticDialect
+ MLIRFuncDialect
+ MLIRIR
+ MLIRPass
+ MLIRTransforms
+ iree::compiler::Dialect::HAL::Conversion::HALToVM
+ iree::compiler::Dialect::Modules::HAL::Loader::IR
+ iree::compiler::Dialect::Util::IR
+ iree::compiler::Dialect::VM::Conversion
+ iree::compiler::Dialect::VM::Conversion::StandardToVM
+ iree::compiler::Dialect::VM::IR
+ PUBLIC
+)
+
+### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ###
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/HALLoaderToVM/ConvertHALLoaderToVM.cpp b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/HALLoaderToVM/ConvertHALLoaderToVM.cpp
new file mode 100644
index 0000000..5fc387f
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/HALLoaderToVM/ConvertHALLoaderToVM.cpp
@@ -0,0 +1,146 @@
+// Copyright 2022 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 "iree/compiler/Dialect/Modules/HAL/Loader/Conversion/HALLoaderToVM/ConvertHALLoaderToVM.h"
+
+#include "iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertHALToVM.h"
+#include "iree/compiler/Dialect/Util/IR/UtilTypes.h"
+#include "iree/compiler/Dialect/VM/Conversion/ConversionTarget.h"
+#include "iree/compiler/Dialect/VM/Conversion/ImportUtils.h"
+#include "iree/compiler/Dialect/VM/Conversion/TypeConverter.h"
+#include "iree/compiler/Dialect/VM/IR/VMOps.h"
+#include "mlir/IR/Attributes.h"
+#include "mlir/IR/Builders.h"
+#include "mlir/IR/BuiltinOps.h"
+#include "mlir/IR/BuiltinTypes.h"
+#include "mlir/IR/Matchers.h"
+#include "mlir/IR/SymbolTable.h"
+#include "mlir/Transforms/DialectConversion.h"
+
+namespace mlir {
+namespace iree_compiler {
+
+namespace {
+
+// Casts |value| to i32 if it is not already.
+static Value castToI32(Value value, OpBuilder &builder) {
+ if (value.getType().isInteger(32)) return value;
+ return builder.createOrFold<IREE::VM::TruncI64I32Op>(
+ value.getLoc(), builder.getI32Type(), value);
+}
+
+struct ExecutableLoadOpConversion
+ : public OpConversionPattern<IREE::HAL::Loader::ExecutableLoadOp> {
+ ExecutableLoadOpConversion(MLIRContext *context, SymbolTable &importSymbols,
+ TypeConverter &typeConverter, StringRef importName)
+ : OpConversionPattern(context) {
+ importOp = importSymbols.lookup<IREE::VM::ImportOp>(importName);
+ assert(importOp);
+ }
+ LogicalResult matchAndRewrite(
+ IREE::HAL::Loader::ExecutableLoadOp loadOp, OpAdaptor adaptor,
+ ConversionPatternRewriter &rewriter) const override {
+ // Get format string as a rodata blob.
+ auto executableFormatStr = rewriter.create<IREE::VM::RodataInlineOp>(
+ loadOp.getLoc(), loadOp.getFormatAttr());
+
+ // Pack constants, if any.
+ auto constantBuffer = createPackedConstantBuffer(
+ loadOp.getLoc(), adaptor.getConstants(), rewriter);
+
+ auto importType = importOp.getFunctionType();
+ auto callOp = rewriter.replaceOpWithNewOp<IREE::VM::CallOp>(
+ loadOp, SymbolRefAttr::get(importOp), importType.getResults(),
+ ValueRange{
+ executableFormatStr,
+ adaptor.getData(),
+ constantBuffer,
+ });
+ copyImportAttrs(importOp, callOp);
+
+ return success();
+ }
+
+ private:
+ mutable IREE::VM::ImportOp importOp;
+};
+
+struct ExecutableDispatchOpConversion
+ : public OpConversionPattern<IREE::HAL::Loader::ExecutableDispatchOp> {
+ ExecutableDispatchOpConversion(MLIRContext *context,
+ SymbolTable &importSymbols,
+ TypeConverter &typeConverter,
+ StringRef importName)
+ : OpConversionPattern(context) {
+ importOp = importSymbols.lookup<IREE::VM::ImportOp>(importName);
+ assert(importOp);
+ }
+ LogicalResult matchAndRewrite(
+ IREE::HAL::Loader::ExecutableDispatchOp dispatchOp, OpAdaptor adaptor,
+ ConversionPatternRewriter &rewriter) const override {
+ auto entryPoint = rewriter.create<IREE::VM::ConstI32Op>(
+ dispatchOp.getLoc(),
+ static_cast<int32_t>(adaptor.getEntryPoint().getZExtValue()));
+ SmallVector<Value, 8> callOperands = {
+ adaptor.getExecutable(),
+ entryPoint,
+ castToI32(adaptor.getWorkgroupX(), rewriter),
+ castToI32(adaptor.getWorkgroupY(), rewriter),
+ castToI32(adaptor.getWorkgroupZ(), rewriter),
+ };
+ auto pushConstants = adaptor.getPushConstants();
+ SmallVector<int16_t, 5> segmentSizes = {
+ /*executable=*/-1,
+ /*entry_point=*/-1,
+ /*workgroup_x=*/-1,
+ /*workgroup_y=*/-1,
+ /*workgroup_z=*/-1,
+ /*push_constants=*/
+ static_cast<int16_t>(pushConstants.size()),
+ /*bindings=*/
+ static_cast<int16_t>(adaptor.getBindingBuffers().size()),
+ };
+ callOperands.append(pushConstants.begin(), pushConstants.end());
+ for (auto it :
+ llvm::zip(adaptor.getBindingBuffers(), adaptor.getBindingOffsets(),
+ adaptor.getBindingLengths())) {
+ callOperands.push_back(std::get<0>(it));
+ callOperands.push_back(
+ castToImportType(std::get<1>(it), rewriter.getI64Type(), rewriter));
+ callOperands.push_back(
+ castToImportType(std::get<2>(it), rewriter.getI64Type(), rewriter));
+ }
+ auto importType = importOp.getFunctionType();
+ auto callOp = rewriter.replaceOpWithNewOp<IREE::VM::CallVariadicOp>(
+ dispatchOp, SymbolRefAttr::get(importOp), importType.getResults(),
+ segmentSizes, importType.getInputs(), callOperands);
+ copyImportAttrs(importOp, callOp);
+ return success();
+ }
+
+ private:
+ mutable IREE::VM::ImportOp importOp;
+};
+
+} // namespace
+
+void populateHALLoaderToVMPatterns(MLIRContext *context,
+ ConversionTarget &conversionTarget,
+ TypeConverter &typeConverter,
+ SymbolTable &importSymbols,
+ RewritePatternSet &patterns) {
+ patterns.insert<
+ VMImportOpConversion<IREE::HAL::Loader::ExecutableQuerySupportOp>>(
+ context, importSymbols, typeConverter,
+ "hal_loader.executable.query_support");
+ patterns.insert<ExecutableLoadOpConversion>(
+ context, importSymbols, typeConverter, "hal_loader.executable.load");
+ patterns.insert<ExecutableDispatchOpConversion>(
+ context, importSymbols, typeConverter, "hal_loader.executable.dispatch");
+}
+
+} // namespace iree_compiler
+} // namespace mlir
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/HALLoaderToVM/ConvertHALLoaderToVM.h b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/HALLoaderToVM/ConvertHALLoaderToVM.h
new file mode 100644
index 0000000..9a892d5
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/HALLoaderToVM/ConvertHALLoaderToVM.h
@@ -0,0 +1,27 @@
+// Copyright 2022 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
+
+#ifndef IREE_COMPILER_DIALECT_MODULES_HAL_LOADER_CONVERSION_HALLOADER_CONVERTHALLOADERTOVM_H_
+#define IREE_COMPILER_DIALECT_MODULES_HAL_LOADER_CONVERSION_HALLOADER_CONVERTHALLOADERTOVM_H_
+
+#include "iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderOps.h"
+#include "mlir/Pass/Pass.h"
+#include "mlir/Transforms/DialectConversion.h"
+
+namespace mlir {
+namespace iree_compiler {
+
+// Populates conversion patterns from the hal_loader dialect to the VM dialect.
+void populateHALLoaderToVMPatterns(MLIRContext *context,
+ ConversionTarget &conversionTarget,
+ TypeConverter &typeConverter,
+ SymbolTable &importSymbols,
+ RewritePatternSet &patterns);
+
+} // namespace iree_compiler
+} // namespace mlir
+
+#endif // IREE_COMPILER_DIALECT_MODULES_HAL_LOADER_CONVERSION_HALLOADER_CONVERTHALLOADERTOVM_H_
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/HALLoaderToVM/test/BUILD b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/HALLoaderToVM/test/BUILD
new file mode 100644
index 0000000..4c13683
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/HALLoaderToVM/test/BUILD
@@ -0,0 +1,28 @@
+# Copyright 2022 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
+
+load("//build_tools/bazel:iree_lit_test.bzl", "iree_lit_test_suite")
+load("//build_tools/bazel:enforce_glob.bzl", "enforce_glob")
+
+package(
+ features = ["layering_check"],
+ licenses = ["notice"], # Apache 2.0
+)
+
+iree_lit_test_suite(
+ name = "lit",
+ srcs = enforce_glob(
+ [
+ "executable_ops.mlir",
+ ],
+ include = ["*.mlir"],
+ ),
+ cfg = "//compiler:lit.cfg.py",
+ tools = [
+ "//tools:iree-opt",
+ "@llvm-project//llvm:FileCheck",
+ ],
+)
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/HALLoaderToVM/test/CMakeLists.txt b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/HALLoaderToVM/test/CMakeLists.txt
new file mode 100644
index 0000000..dbfca8a
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/HALLoaderToVM/test/CMakeLists.txt
@@ -0,0 +1,23 @@
+################################################################################
+# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from #
+# compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/HALLoaderToVM/test/BUILD#
+# #
+# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary #
+# CMake-only content. #
+# #
+# To disable autogeneration for this file entirely, delete this header. #
+################################################################################
+
+iree_add_all_subdirs()
+
+iree_lit_test_suite(
+ NAME
+ lit
+ SRCS
+ "executable_ops.mlir"
+ TOOLS
+ FileCheck
+ iree-opt
+)
+
+### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ###
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/HALLoaderToVM/test/executable_ops.mlir b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/HALLoaderToVM/test/executable_ops.mlir
new file mode 100644
index 0000000..8328041
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/HALLoaderToVM/test/executable_ops.mlir
@@ -0,0 +1,54 @@
+// RUN: iree-opt --iree-vm-target-index-bits=64 --split-input-file \
+// RUN: --iree-vm-conversion --canonicalize %s | FileCheck %s
+
+// CHECK-LABEL: @executableLoad
+// CHECK-SAME: (%[[EXECUTABLE_DATA:.+]]: !vm.buffer)
+func.func @executableLoad(%executable_data: !util.buffer) -> !hal.executable {
+ // CHECK-DAG: %[[CONSTANTS:.+]] = vm.const.ref.zero : !vm.buffer
+ // CHECK-DAG: %[[FORMAT_STR:.+]] = vm.rodata.inline {{.+}} : !vm.buffer = "executable_format"
+ // CHECK: %[[EXECUTABLE:.+]] = vm.call @hal_loader.executable.load(%[[FORMAT_STR]], %[[EXECUTABLE_DATA]], %[[CONSTANTS]])
+ %executable = hal_loader.executable.load format("executable_format") data(%executable_data) : !hal.executable
+ // CHECK: return %[[EXECUTABLE]]
+ return %executable : !hal.executable
+}
+
+// -----
+
+// CHECK-LABEL: @executableDispatch
+// CHECK-SAME: (%[[EXECUTABLE:.+]]: !vm.ref<!hal.executable>,
+// CHECK-SAME: %[[BUFFER0:.+]]: !vm.buffer, %[[BUFFER1:.+]]: !vm.buffer)
+func.func @executableDispatch(%executable: !hal.executable, %buffer0: !util.buffer, %buffer1: !util.buffer) {
+ // CHECK-DAG: %[[COUNT_X:.+]] = vm.const.i32 1000
+ %count_x = arith.constant 1000 : index
+ // CHECK-DAG: %[[COUNT_Y:.+]] = vm.const.i32 1001
+ %count_y = arith.constant 1001 : index
+ // CHECK-DAG: %[[COUNT_Z:.+]] = vm.const.i32 1002
+ %count_z = arith.constant 1002 : index
+ // CHECK-DAG: %[[CONSTANT0:.+]] = vm.const.i32 4
+ %constant0 = arith.constant 4 : i32
+ // CHECK-DAG: %[[CONSTANT1:.+]] = vm.const.i32 5
+ %constant1 = arith.constant 5 : i32
+ // CHECK-DAG: %[[OFFSET0:.+]] = vm.const.i64 100
+ %offset0 = arith.constant 100 : index
+ // CHECK-DAG: %[[LENGTH0:.+]] = vm.const.i64 128
+ %length0 = arith.constant 128 : index
+ // CHECK-DAG: %[[OFFSET1:.+]] = vm.const.i64 200
+ %offset1 = arith.constant 200 : index
+ // CHECK-DAG: %[[LENGTH1:.+]] = vm.const.i64 256
+ %length1 = arith.constant 256 : index
+ // CHECK: vm.call.variadic @hal_loader.executable.dispatch
+ hal_loader.executable.dispatch
+ // CHECK-SAME: %[[EXECUTABLE]], %c16
+ executable(%executable : !hal.executable)[16]
+ // CHECK-SAME: %[[COUNT_X]], %[[COUNT_Y]], %[[COUNT_Z]]
+ workgroups([%count_x, %count_y, %count_z])
+ // CHECK-SAME: [%[[CONSTANT0]], %[[CONSTANT1]]]
+ constants([%constant0, %constant1])
+ bindings([
+ // CHECK-SAME: (%[[BUFFER0]], %[[OFFSET0]], %[[LENGTH0]])
+ (%buffer0 : !util.buffer)[%offset0, %length0],
+ // CHECK-SAME: (%[[BUFFER1]], %[[OFFSET1]], %[[LENGTH1]])
+ (%buffer1 : !util.buffer)[%offset1, %length1]
+ ])
+ return
+}
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/StreamToHALLoader/BUILD b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/StreamToHALLoader/BUILD
new file mode 100644
index 0000000..2c2b26b
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/StreamToHALLoader/BUILD
@@ -0,0 +1,41 @@
+# Copyright 2022 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
+
+load("//build_tools/bazel:build_defs.oss.bzl", "iree_compiler_cc_library")
+
+package(
+ default_visibility = ["//visibility:public"],
+ features = ["layering_check"],
+ licenses = ["notice"], # Apache 2.0
+)
+
+iree_compiler_cc_library(
+ name = "StreamToHALLoader",
+ srcs = [
+ "ConvertStreamToHALLoader.cpp",
+ ],
+ hdrs = [
+ "ConvertStreamToHALLoader.h",
+ ],
+ deps = [
+ "//compiler/src/iree/compiler/Dialect/HAL/Conversion",
+ "//compiler/src/iree/compiler/Dialect/HAL/IR",
+ "//compiler/src/iree/compiler/Dialect/HAL/IR:HALDialect",
+ "//compiler/src/iree/compiler/Dialect/HAL/Target",
+ "//compiler/src/iree/compiler/Dialect/HAL/Utils",
+ "//compiler/src/iree/compiler/Dialect/Modules/HAL/Inline/IR",
+ "//compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR",
+ "//compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR:HALLoaderDialect",
+ "//compiler/src/iree/compiler/Dialect/Stream/IR",
+ "//compiler/src/iree/compiler/Dialect/Util/IR",
+ "@llvm-project//llvm:Support",
+ "@llvm-project//mlir:ArithmeticDialect",
+ "@llvm-project//mlir:FuncDialect",
+ "@llvm-project//mlir:IR",
+ "@llvm-project//mlir:Pass",
+ "@llvm-project//mlir:Transforms",
+ ],
+)
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/StreamToHALLoader/CMakeLists.txt b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/StreamToHALLoader/CMakeLists.txt
new file mode 100644
index 0000000..9eee0b6
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/StreamToHALLoader/CMakeLists.txt
@@ -0,0 +1,40 @@
+################################################################################
+# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from #
+# compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/StreamToHALLoader/BUILD#
+# #
+# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary #
+# CMake-only content. #
+# #
+# To disable autogeneration for this file entirely, delete this header. #
+################################################################################
+
+iree_add_all_subdirs()
+
+iree_cc_library(
+ NAME
+ StreamToHALLoader
+ HDRS
+ "ConvertStreamToHALLoader.h"
+ SRCS
+ "ConvertStreamToHALLoader.cpp"
+ DEPS
+ LLVMSupport
+ MLIRArithmeticDialect
+ MLIRFuncDialect
+ MLIRIR
+ MLIRPass
+ MLIRTransforms
+ iree::compiler::Dialect::HAL::Conversion
+ iree::compiler::Dialect::HAL::IR
+ iree::compiler::Dialect::HAL::IR::HALDialect
+ iree::compiler::Dialect::HAL::Target
+ iree::compiler::Dialect::HAL::Utils
+ iree::compiler::Dialect::Modules::HAL::Inline::IR
+ iree::compiler::Dialect::Modules::HAL::Loader::IR
+ iree::compiler::Dialect::Modules::HAL::Loader::IR::HALLoaderDialect
+ iree::compiler::Dialect::Stream::IR
+ iree::compiler::Dialect::Util::IR
+ PUBLIC
+)
+
+### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ###
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/StreamToHALLoader/ConvertStreamToHALLoader.cpp b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/StreamToHALLoader/ConvertStreamToHALLoader.cpp
new file mode 100644
index 0000000..a4ca7ba
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/StreamToHALLoader/ConvertStreamToHALLoader.cpp
@@ -0,0 +1,150 @@
+// Copyright 2022 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 "iree/compiler/Dialect/Modules/HAL/Loader/Conversion/StreamToHALLoader/ConvertStreamToHALLoader.h"
+
+#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
+#include "iree/compiler/Dialect/HAL/IR/HALTypes.h"
+#include "iree/compiler/Dialect/Modules/HAL/Inline/IR/HALInlineOps.h"
+#include "iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderDialect.h"
+#include "iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderOps.h"
+#include "iree/compiler/Dialect/Stream/IR/StreamDialect.h"
+#include "iree/compiler/Dialect/Stream/IR/StreamOps.h"
+#include "iree/compiler/Dialect/Stream/IR/StreamTypes.h"
+#include "iree/compiler/Dialect/Util/IR/UtilOps.h"
+#include "mlir/Dialect/Arithmetic/IR/Arithmetic.h"
+#include "mlir/Dialect/Func/IR/FuncOps.h"
+#include "mlir/Transforms/DialectConversion.h"
+
+namespace mlir {
+namespace iree_compiler {
+
+namespace {
+
+// Returns the !util.buffer from the given converted resource, which may be
+// either a !util.buffer or an external !hal.buffer.
+static Value getResourceBuffer(Location loc, Value resource,
+ OpBuilder &builder) {
+ if (resource.getType().isa<IREE::HAL::BufferType>()) {
+ // Get the storage of the buffer; the returned buffer is already a subspan.
+ return builder.createOrFold<IREE::HAL::Inline::BufferStorageOp>(loc,
+ resource);
+ }
+ return resource;
+}
+
+// Converts a dispatch command into an inline executable dispatch.
+struct CmdDispatchOpPattern
+ : public OpConversionPattern<IREE::Stream::CmdDispatchOp> {
+ CmdDispatchOpPattern(TypeConverter &typeConverter, MLIRContext *context)
+ : OpConversionPattern(typeConverter, context, PatternBenefit(10000)) {}
+ LogicalResult matchAndRewrite(
+ IREE::Stream::CmdDispatchOp dispatchOp, OpAdaptor adaptor,
+ ConversionPatternRewriter &rewriter) const override {
+ auto loc = dispatchOp.getLoc();
+
+ // Get the handle to the executable that is compatible with our device.
+ auto executableOp =
+ cast<IREE::HAL::ExecutableOp>(SymbolTable::lookupNearestSymbolFrom(
+ dispatchOp, dispatchOp.getEntryPoint().getRootReference()));
+ assert(executableOp && "dispatch target executable op not found");
+
+ // For now we aren't doing loader support checks. We should, though.
+ auto variantOps = executableOp.getOps<IREE::HAL::ExecutableVariantOp>();
+ if (std::distance(variantOps.begin(), variantOps.end()) > 1) {
+ return rewriter.notifyMatchFailure(dispatchOp,
+ "only one variant is supported today");
+ }
+
+ // Lookup executable reference.
+ auto lookupOp = rewriter.create<IREE::HAL::Loader::ExecutableLookupOp>(
+ loc, rewriter.getType<IREE::HAL::ExecutableType>(),
+ executableOp.getName());
+
+ // TODO(benvanik): a real switch op. For now we inline what the
+ // hal.device.switch op does.
+ for (auto variantOp : variantOps) {
+ auto exportOps = variantOp.getOps<IREE::HAL::ExecutableExportOp>();
+ auto exportIt =
+ llvm::find_if(exportOps, [&](IREE::HAL::ExecutableExportOp op) {
+ return op.getNameAttr() ==
+ dispatchOp.getEntryPoint().getLeafReference();
+ });
+ if (exportIt == exportOps.end()) {
+ return variantOp.emitError()
+ << "hal.executable.variant is missing the entry point for "
+ << dispatchOp.getEntryPoint();
+ }
+ auto exportOp = *exportIt;
+
+ // TODO(benvanik): check variant target:
+ // if (variantOp.target().getMatchExpression()) { dispatch }
+ dispatchVariant(dispatchOp, adaptor, executableOp, variantOp, exportOp,
+ lookupOp.getResult(), rewriter);
+ }
+
+ rewriter.eraseOp(dispatchOp);
+ return success();
+ }
+
+ void dispatchVariant(IREE::Stream::CmdDispatchOp dispatchOp,
+ OpAdaptor adaptor, IREE::HAL::ExecutableOp executableOp,
+ IREE::HAL::ExecutableVariantOp variantOp,
+ IREE::HAL::ExecutableExportOp exportOp, Value executable,
+ OpBuilder &builder) const {
+ auto loc = dispatchOp.getLoc();
+
+ // Push constant values.
+ // TODO(#5322): symbolic push constant names on the hal.interface so we can
+ // sparsely pack these.
+ SmallVector<Value> pushConstants;
+ for (auto operand : adaptor.getUniformOperands()) {
+ assert(operand.getType().isInteger(32) &&
+ "expected only i32 values after iree-hal-pack-dispatch-operands");
+ pushConstants.push_back(operand);
+ }
+
+ // Push descriptor bindings.
+ SmallVector<Value> bindingBuffers;
+ SmallVector<Value> bindingOffsets;
+ SmallVector<Value> bindingLengths;
+ for (unsigned i = 0; i < adaptor.getResources().size(); ++i) {
+ auto buffer = getResourceBuffer(loc, adaptor.getResources()[i], builder);
+ bindingBuffers.push_back(buffer);
+ bindingOffsets.push_back(adaptor.getResourceOffsets()[i]);
+ bindingLengths.push_back(adaptor.getResourceLengths()[i]);
+ }
+
+ // Dispatch with a target-specific workgroup count.
+ auto exportSymRef =
+ SymbolRefAttr::get(builder.getContext(), executableOp.getName(),
+ {SymbolRefAttr::get(exportOp->getParentOp()),
+ SymbolRefAttr::get(exportOp)});
+ auto workgroupCount = exportOp.calculateWorkgroupCount(
+ loc, /*device=*/nullptr, adaptor.getWorkload(), builder);
+ builder.create<IREE::HAL::Loader::ExecutableDispatchSymbolOp>(
+ loc, executable, exportSymRef, workgroupCount[0], workgroupCount[1],
+ workgroupCount[2], pushConstants, bindingBuffers, bindingOffsets,
+ bindingLengths);
+ }
+};
+
+} // namespace
+
+void populateStreamToHALLoaderPatterns(MLIRContext *context,
+ ConversionTarget &conversionTarget,
+ TypeConverter &typeConverter,
+ RewritePatternSet &patterns) {
+ // Executables are taken care of after serialization by the
+ // MaterializeExecutables pass. We allow them to pass through for now.
+ conversionTarget.addLegalOp<IREE::HAL::ExecutableOp>();
+ conversionTarget.markOpRecursivelyLegal<IREE::HAL::ExecutableOp>();
+
+ patterns.insert<CmdDispatchOpPattern>(typeConverter, context);
+}
+
+} // namespace iree_compiler
+} // namespace mlir
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/StreamToHALLoader/ConvertStreamToHALLoader.h b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/StreamToHALLoader/ConvertStreamToHALLoader.h
new file mode 100644
index 0000000..203bdc0
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/StreamToHALLoader/ConvertStreamToHALLoader.h
@@ -0,0 +1,25 @@
+// Copyright 2022 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
+
+#ifndef IREE_COMPILER_DIALECT_MODULES_HAL_LOADER_CONVERSION_STREAMTOHALLOADER_CONVERTSTREAMTOHALLOADER_H_
+#define IREE_COMPILER_DIALECT_MODULES_HAL_LOADER_CONVERSION_STREAMTOHALLOADER_CONVERTSTREAMTOHALLOADER_H_
+
+#include "mlir/Pass/Pass.h"
+#include "mlir/Transforms/DialectConversion.h"
+
+namespace mlir {
+namespace iree_compiler {
+
+// Populates conversion patterns for stream->HAL (loader).
+void populateStreamToHALLoaderPatterns(MLIRContext *context,
+ ConversionTarget &conversionTarget,
+ TypeConverter &typeConverter,
+ RewritePatternSet &patterns);
+
+} // namespace iree_compiler
+} // namespace mlir
+
+#endif // IREE_COMPILER_DIALECT_MODULES_HAL_LOADER_CONVERSION_STREAMTOHALLOADER_CONVERTSTREAMTOHALLOADER_H_
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/StreamToHALLoader/test/BUILD b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/StreamToHALLoader/test/BUILD
new file mode 100644
index 0000000..e3365e1
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/StreamToHALLoader/test/BUILD
@@ -0,0 +1,28 @@
+# Copyright 2021 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
+
+load("//build_tools/bazel:iree_lit_test.bzl", "iree_lit_test_suite")
+load("//build_tools/bazel:enforce_glob.bzl", "enforce_glob")
+
+package(
+ features = ["layering_check"],
+ licenses = ["notice"], # Apache 2.0
+)
+
+iree_lit_test_suite(
+ name = "lit",
+ srcs = enforce_glob(
+ [
+ "cmd_ops.mlir",
+ ],
+ include = ["*.mlir"],
+ ),
+ cfg = "//compiler:lit.cfg.py",
+ tools = [
+ "//tools:iree-opt",
+ "@llvm-project//llvm:FileCheck",
+ ],
+)
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/StreamToHALLoader/test/CMakeLists.txt b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/StreamToHALLoader/test/CMakeLists.txt
new file mode 100644
index 0000000..a3c0aba
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/StreamToHALLoader/test/CMakeLists.txt
@@ -0,0 +1,23 @@
+################################################################################
+# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from #
+# compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/StreamToHALLoader/test/BUILD#
+# #
+# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary #
+# CMake-only content. #
+# #
+# To disable autogeneration for this file entirely, delete this header. #
+################################################################################
+
+iree_add_all_subdirs()
+
+iree_lit_test_suite(
+ NAME
+ lit
+ SRCS
+ "cmd_ops.mlir"
+ TOOLS
+ FileCheck
+ iree-opt
+)
+
+### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ###
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/StreamToHALLoader/test/cmd_ops.mlir b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/StreamToHALLoader/test/cmd_ops.mlir
new file mode 100644
index 0000000..1008469
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/StreamToHALLoader/test/cmd_ops.mlir
@@ -0,0 +1,91 @@
+// RUN: iree-opt --split-input-file --iree-hal-loader-conversion %s | FileCheck %s
+
+// NOTE: all other stream.cmd.* ops are handled by the hal_inline conversions.
+
+// Executables are required to translate the dispatch calls.
+#executable_layout = #hal.executable.layout<push_constants = 0, sets = [
+ #hal.descriptor_set.layout<0, bindings = [
+ #hal.descriptor_set.binding<4, storage_buffer>
+ ]>,
+ #hal.descriptor_set.layout<1, bindings = [
+ #hal.descriptor_set.binding<5, storage_buffer>
+ ]>
+]>
+hal.executable private @ex {
+ hal.executable.variant public @variant, target = #hal.executable.target<"llvm", "embedded-elf-x86_64"> {
+ hal.executable.export public @dispatch ordinal(16) layout(#executable_layout) {
+ ^bb0(%device: !hal.device, %workload_x: index, %workload_y: index):
+ %count_x = affine.apply affine_map<()[s0] -> (s0 ceildiv 4)>()[%workload_x]
+ %count_y = affine.apply affine_map<()[s0] -> (s0 ceildiv 4)>()[%workload_y]
+ %count_z = arith.constant 1 : index
+ hal.return %count_x, %count_y, %count_z : index, index, index
+ }
+ builtin.module {
+ // Opaque at this point (in some target-specific dialects).
+ }
+ }
+}
+
+// NOTE: %buffer0 is transient and will map to a !util.buffer, while
+// %buffer1 is external and will map to a !hal.buffer.
+
+// CHECK-LABEL: @cmdDispatch
+// CHECK-SAME: (%[[BUFFER0:.+]]: !util.buffer, %[[BUFFER0_SIZE:.+]]: index,
+// CHECK-SAME: %[[BUFFER1:.+]]: !hal.buffer, %[[BUFFER1_SIZE:.+]]: index)
+func.func @cmdDispatch(%buffer0: !stream.resource<transient>, %buffer0_size: index,
+ %buffer1: !stream.resource<external>, %buffer1_size: index) -> !stream.timepoint {
+ // (ends up by the dispatch below)
+ %workload_x = arith.constant 1000 : index
+ %workload_y = arith.constant 1001 : index
+
+ // CHECK-DAG: %[[CONSTANT0:.+]] = arith.constant 4
+ %constant0 = arith.constant 4 : i32
+ // CHECK-DAG: %[[CONSTANT1:.+]] = arith.constant 5
+ %constant1 = arith.constant 5 : i32
+
+ // CHECK: %[[BUFFER0_REL_OFFSET:.+]] = arith.constant 200
+ %buffer0_offset = arith.constant 200 : index
+ // CHECK: %[[BUFFER0_REL_LENGTH:.+]] = arith.constant 128
+ %buffer0_length = arith.constant 128 : index
+ // CHECK: %[[BUFFER1_REL_OFFSET:.+]] = arith.constant 300
+ %buffer1_offset = arith.constant 300 : index
+ // CHECK: %[[BUFFER1_REL_LENGTH:.+]] = arith.constant 256
+ %buffer1_length = arith.constant 256 : index
+
+ %fence = stream.cmd.execute with(%buffer0 as %buffer0_inner: !stream.resource<transient>{%buffer0_size},
+ %buffer1 as %buffer1_inner: !stream.resource<external>{%buffer1_size}) {
+ // Lookup the loaded executable (resolved during iree-hal-loader-materialize-executables):
+ // CHECK: %[[EXECUTABLE:.+]] = hal_loader.executable.lookup executable(@ex) : !hal.executable
+
+ // %buffer1 is external and a subspan is needed to resolve the absolute
+ // storage range. This will (mostly) eventually fold/canonicalize away.
+ // CHECK: %[[BUFFER1_STORAGE:.+]] = hal_inline.buffer.storage<%[[BUFFER1]]
+
+ // Workload calculation gets inlined and folds during conversion; this is
+ // the original worload ceildiv 4 on x/y:
+ // CHECK-DAG: %[[COUNT_X:.+]] = arith.constant 250
+ // CHECK-DAG: %[[COUNT_Y:.+]] = arith.constant 251
+ // CHECK-DAG: %[[COUNT_Z:.+]] = arith.constant 1
+
+ // CHECK: hal_loader.executable.dispatch
+ // CHECK-SAME: executable(%[[EXECUTABLE]] : !hal.executable)
+ // CHECK-SAME: target(@ex::@variant::@dispatch)
+ // CHECK-SAME: workgroups([%[[COUNT_X]], %[[COUNT_Y]], %[[COUNT_Z]]])
+ // CHECK-SAME: constants([%[[CONSTANT0]], %[[CONSTANT1]]])
+ // CHECK-SAME: bindings([
+ // CHECK-NEXT: (%[[BUFFER0]] : !util.buffer)[%[[BUFFER0_REL_OFFSET]], %[[BUFFER0_REL_LENGTH]]],
+ // CHECK-NEXT: (%[[BUFFER1_STORAGE]] : !util.buffer)[%[[BUFFER1_REL_OFFSET]], %[[BUFFER1_REL_LENGTH]]]
+ // CHECK-NEXT: ])
+ stream.cmd.dispatch @ex::@dispatch[%workload_x, %workload_y](%constant0, %constant1 : i32, i32) {
+ ro %buffer0_inner[%buffer0_offset for %buffer0_length] : !stream.resource<transient>{%buffer0_size},
+ wo %buffer1_inner[%buffer1_offset for %buffer1_length] : !stream.resource<external>{%buffer1_size}
+ } attributes {
+ hal.interface.bindings = [
+ #hal.interface.binding<0, 4>,
+ #hal.interface.binding<1, 5>
+ ]
+ }
+ } => !stream.timepoint
+ // CHECK: return %c0
+ return %fence : !stream.timepoint
+}
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/BUILD b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/BUILD
new file mode 100644
index 0000000..732243e
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/BUILD
@@ -0,0 +1,114 @@
+# Copyright 2022 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
+
+load("@llvm-project//mlir:tblgen.bzl", "td_library")
+load("//build_tools/bazel:build_defs.oss.bzl", "iree_compiler_cc_library")
+load("//build_tools/bazel:enforce_glob.bzl", "enforce_glob")
+load("//build_tools/bazel:iree_tablegen.bzl", "iree_gentbl_cc_library", "iree_tablegen_doc")
+
+package(
+ default_visibility = ["//visibility:public"],
+ features = ["layering_check"],
+ licenses = ["notice"], # Apache 2.0
+)
+
+exports_files(["HALLoaderOps.td"])
+
+td_library(
+ name = "td_files",
+ srcs = enforce_glob(
+ [
+ "HALLoaderBase.td",
+ "HALLoaderOps.td",
+ ],
+ include = ["*.td"],
+ ),
+ deps = [
+ "//compiler/src/iree/compiler/Dialect/HAL/IR:td_files",
+ "//compiler/src/iree/compiler/Dialect/Util/IR:td_files",
+ "@llvm-project//mlir:FuncTdFiles",
+ "@llvm-project//mlir:OpBaseTdFiles",
+ ],
+)
+
+iree_compiler_cc_library(
+ name = "IR",
+ srcs = [
+ "HALLoaderOps.cpp",
+ ],
+ hdrs = [
+ "HALLoaderOps.h",
+ "HALLoaderOps.h.inc",
+ ],
+ textual_hdrs = [
+ "HALLoaderOps.cpp.inc",
+ ],
+ deps = [
+ ":HALLoaderOpsGen",
+ "//compiler/src/iree/compiler/Dialect/HAL/IR",
+ "//compiler/src/iree/compiler/Dialect/Util/IR",
+ "//compiler/src/iree/compiler/Dialect/VM/IR",
+ "@llvm-project//llvm:Support",
+ "@llvm-project//mlir:ArithmeticDialect",
+ "@llvm-project//mlir:FuncDialect",
+ "@llvm-project//mlir:IR",
+ "@llvm-project//mlir:SideEffectInterfaces",
+ "@llvm-project//mlir:Support",
+ "@llvm-project//mlir:TransformUtils",
+ "@llvm-project//mlir:TranslateLib",
+ ],
+)
+
+iree_compiler_cc_library(
+ name = "HALLoaderDialect",
+ srcs = ["HALLoaderDialect.cpp"],
+ hdrs = ["HALLoaderDialect.h"],
+ deps = [
+ ":IR",
+ "//compiler/src/iree/compiler/Dialect/Modules/HAL/Loader:hal_loader_imports",
+ "//compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/HALLoaderToVM",
+ "//compiler/src/iree/compiler/Dialect/VM/Conversion",
+ "@llvm-project//llvm:Support",
+ "@llvm-project//mlir:FuncDialect",
+ "@llvm-project//mlir:IR",
+ "@llvm-project//mlir:Parser",
+ "@llvm-project//mlir:Support",
+ "@llvm-project//mlir:TransformUtils",
+ ],
+)
+
+iree_gentbl_cc_library(
+ name = "HALLoaderOpsGen",
+ tbl_outs = [
+ (
+ ["--gen-op-decls"],
+ "HALLoaderOps.h.inc",
+ ),
+ (
+ ["--gen-op-defs"],
+ "HALLoaderOps.cpp.inc",
+ ),
+ ],
+ tblgen = "@llvm-project//mlir:mlir-tblgen",
+ td_file = "HALLoaderOps.td",
+ deps = [":td_files"],
+)
+
+iree_tablegen_doc(
+ name = "HALLoaderDialecDocGen",
+ tbl_outs = [
+ (
+ [
+ "--dialect=hal_loader",
+ "--gen-dialect-doc",
+ ],
+ "HALLoaderDialect.md",
+ ),
+ ],
+ tblgen = "@llvm-project//mlir:mlir-tblgen",
+ td_file = "HALLoaderOps.td",
+ deps = [":td_files"],
+)
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/CMakeLists.txt b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/CMakeLists.txt
new file mode 100644
index 0000000..d673a27
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/CMakeLists.txt
@@ -0,0 +1,79 @@
+################################################################################
+# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from #
+# compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/BUILD #
+# #
+# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary #
+# CMake-only content. #
+# #
+# To disable autogeneration for this file entirely, delete this header. #
+################################################################################
+
+iree_add_all_subdirs()
+
+iree_cc_library(
+ NAME
+ IR
+ HDRS
+ "HALLoaderOps.h"
+ "HALLoaderOps.h.inc"
+ TEXTUAL_HDRS
+ "HALLoaderOps.cpp.inc"
+ SRCS
+ "HALLoaderOps.cpp"
+ DEPS
+ ::HALLoaderOpsGen
+ LLVMSupport
+ MLIRArithmeticDialect
+ MLIRFuncDialect
+ MLIRIR
+ MLIRSideEffectInterfaces
+ MLIRSupport
+ MLIRTransformUtils
+ MLIRTranslateLib
+ iree::compiler::Dialect::HAL::IR
+ iree::compiler::Dialect::Util::IR
+ iree::compiler::Dialect::VM::IR
+ PUBLIC
+)
+
+iree_cc_library(
+ NAME
+ HALLoaderDialect
+ HDRS
+ "HALLoaderDialect.h"
+ SRCS
+ "HALLoaderDialect.cpp"
+ DEPS
+ ::IR
+ LLVMSupport
+ MLIRFuncDialect
+ MLIRIR
+ MLIRParser
+ MLIRSupport
+ MLIRTransformUtils
+ iree::compiler::Dialect::Modules::HAL::Loader::Conversion::HALLoaderToVM
+ iree::compiler::Dialect::Modules::HAL::Loader::hal_loader_imports
+ iree::compiler::Dialect::VM::Conversion
+ PUBLIC
+)
+
+iree_tablegen_library(
+ NAME
+ HALLoaderOpsGen
+ TD_FILE
+ "HALLoaderOps.td"
+ OUTS
+ --gen-op-decls HALLoaderOps.h.inc
+ --gen-op-defs HALLoaderOps.cpp.inc
+)
+
+iree_tablegen_doc(
+ NAME
+ HALLoaderDialecDocGen
+ TD_FILE
+ "HALLoaderOps.td"
+ OUTS
+ --dialect=hal_loader --gen-dialect-doc HALLoaderDialect.md
+)
+
+### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ###
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderBase.td b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderBase.td
new file mode 100644
index 0000000..13a4f93
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderBase.td
@@ -0,0 +1,56 @@
+// Copyright 2022 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
+
+#ifndef IREE_DIALECT_MODULES_HAL_LOADER_BASE
+#define IREE_DIALECT_MODULES_HAL_LOADER_BASE
+
+include "iree/compiler/Dialect/Util/IR/UtilBase.td"
+
+//===----------------------------------------------------------------------===//
+// IREE HAL inline executable loader dialect
+//===----------------------------------------------------------------------===//
+
+def HALLoader_Dialect : Dialect {
+ let name = "hal_loader";
+ let cppNamespace = "::mlir::iree_compiler::IREE::HAL::Loader";
+ let emitAccessorPrefix = kEmitAccessorPrefix_Prefixed;
+
+ let summary = [{
+ IREE HAL inline executable loader runtime module dialect.
+ }];
+ let description = [{
+ Low-level dialect for dynamically loading executables and dispatching work.
+ Only operates synchronously, single-threaded, and on host-local buffers. Use
+ the full HAL for all other cases.
+
+ This dialect can be used alongside the full HAL but is intended for use in
+ conjunction with the `hal_inline` dialect which also carries the same usage
+ restrictions.
+
+ See `hal_loader.imports.mlir` for the full list of exported functions.
+ }];
+}
+
+//===----------------------------------------------------------------------===//
+// HALLoader enums
+//===----------------------------------------------------------------------===//
+
+//===----------------------------------------------------------------------===//
+// HALLoader types
+//===----------------------------------------------------------------------===//
+
+//===----------------------------------------------------------------------===//
+// HALLoader op traits
+//===----------------------------------------------------------------------===//
+
+//===----------------------------------------------------------------------===//
+// Base HALLoader op classes
+//===----------------------------------------------------------------------===//
+
+class HALLoader_Op<string mnemonic, list<Trait> traits = []> :
+ Op<HALLoader_Dialect, mnemonic, !listconcat(traits, [])> {}
+
+#endif // IREE_DIALECT_MODULES_HAL_LOADER_BASE
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderDialect.cpp b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderDialect.cpp
new file mode 100644
index 0000000..08b6418
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderDialect.cpp
@@ -0,0 +1,63 @@
+// Copyright 2022 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 "iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderDialect.h"
+
+#include "iree/compiler/Dialect/Modules/HAL/Loader/Conversion/HALLoaderToVM/ConvertHALLoaderToVM.h"
+#include "iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderOps.h"
+#include "iree/compiler/Dialect/Modules/HAL/Loader/hal_loader.imports.h"
+#include "iree/compiler/Dialect/VM/Conversion/ConversionDialectInterface.h"
+#include "llvm/Support/SourceMgr.h"
+#include "mlir/IR/DialectImplementation.h"
+#include "mlir/IR/OpImplementation.h"
+#include "mlir/Parser/Parser.h"
+
+namespace mlir {
+namespace iree_compiler {
+namespace IREE {
+namespace HAL {
+namespace Loader {
+
+namespace {
+
+class HALLoaderToVMConversionInterface : public VMConversionDialectInterface {
+ public:
+ using VMConversionDialectInterface::VMConversionDialectInterface;
+
+ OwningOpRef<mlir::ModuleOp> parseVMImportModule() const override {
+ return mlir::parseSourceString<mlir::ModuleOp>(
+ StringRef(iree_hal_loader_imports_create()->data,
+ iree_hal_loader_imports_create()->size),
+ getDialect()->getContext());
+ }
+
+ void populateVMConversionPatterns(
+ SymbolTable &importSymbols, RewritePatternSet &patterns,
+ ConversionTarget &conversionTarget,
+ TypeConverter &typeConverter) const override {
+ conversionTarget.addIllegalDialect<IREE::HAL::Loader::HALLoaderDialect>();
+ populateHALLoaderToVMPatterns(getDialect()->getContext(), conversionTarget,
+ typeConverter, importSymbols, patterns);
+ }
+};
+
+} // namespace
+
+HALLoaderDialect::HALLoaderDialect(MLIRContext *context)
+ : Dialect(getDialectNamespace(), context, TypeID::get<HALLoaderDialect>()) {
+ addInterfaces<HALLoaderToVMConversionInterface>();
+
+#define GET_OP_LIST
+ addOperations<
+#include "iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderOps.cpp.inc"
+ >();
+}
+
+} // namespace Loader
+} // namespace HAL
+} // namespace IREE
+} // namespace iree_compiler
+} // namespace mlir
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderDialect.h b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderDialect.h
new file mode 100644
index 0000000..4676294
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderDialect.h
@@ -0,0 +1,31 @@
+// Copyright 2022 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
+
+#ifndef IREE_COMPILER_DIALECT_MODULES_HAL_LOADER_IR_HALLOADERDIALECT_H_
+#define IREE_COMPILER_DIALECT_MODULES_HAL_LOADER_IR_HALLOADERDIALECT_H_
+
+#include "mlir/IR/Dialect.h"
+#include "mlir/IR/OpDefinition.h"
+
+namespace mlir {
+namespace iree_compiler {
+namespace IREE {
+namespace HAL {
+namespace Loader {
+
+class HALLoaderDialect : public Dialect {
+ public:
+ explicit HALLoaderDialect(MLIRContext *context);
+ static StringRef getDialectNamespace() { return "hal_loader"; }
+};
+
+} // namespace Loader
+} // namespace HAL
+} // namespace IREE
+} // namespace iree_compiler
+} // namespace mlir
+
+#endif // IREE_COMPILER_DIALECT_MODULES_HAL_LOADER_IR_HALLOADERDIALECT_H_
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderOps.cpp b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderOps.cpp
new file mode 100644
index 0000000..14586bd
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderOps.cpp
@@ -0,0 +1,212 @@
+// Copyright 2022 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 "iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderOps.h"
+
+#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
+#include "iree/compiler/Dialect/Util/IR/UtilOps.h"
+#include "iree/compiler/Dialect/Util/IR/UtilTypes.h"
+#include "llvm/ADT/STLExtras.h"
+#include "llvm/Support/SMLoc.h"
+#include "mlir/Dialect/Arithmetic/IR/Arithmetic.h"
+#include "mlir/IR/Attributes.h"
+#include "mlir/IR/Builders.h"
+#include "mlir/IR/OpImplementation.h"
+#include "mlir/IR/SymbolTable.h"
+#include "mlir/IR/TypeUtilities.h"
+
+namespace mlir {
+namespace iree_compiler {
+namespace IREE {
+namespace HAL {
+namespace Loader {
+
+//===----------------------------------------------------------------------===//
+// custom<DispatchBindings>($binding_buffers,
+// type($binding_buffers),
+// $binding_offsets,
+// $binding_lengths)
+//===----------------------------------------------------------------------===//
+
+static ParseResult parseDispatchBindings(
+ OpAsmParser &parser,
+ SmallVectorImpl<OpAsmParser::UnresolvedOperand> &buffers,
+ SmallVectorImpl<Type> &bufferTypes,
+ SmallVectorImpl<OpAsmParser::UnresolvedOperand> &bufferOffsets,
+ SmallVectorImpl<OpAsmParser::UnresolvedOperand> &bufferLengths) {
+ do {
+ OpAsmParser::UnresolvedOperand ordinal;
+ OpAsmParser::UnresolvedOperand buffer;
+ Type bufferType;
+ OpAsmParser::UnresolvedOperand bufferOffset;
+ OpAsmParser::UnresolvedOperand bufferLength;
+ if (failed(parser.parseLParen()) || failed(parser.parseOperand(buffer)) ||
+ failed(parser.parseColonType(bufferType)) ||
+ failed(parser.parseRParen()) || failed(parser.parseLSquare()) ||
+ failed(parser.parseOperand(bufferOffset)) ||
+ failed(parser.parseComma()) ||
+ failed(parser.parseOperand(bufferLength)) ||
+ failed(parser.parseRSquare())) {
+ return failure();
+ }
+ buffers.push_back(buffer);
+ bufferTypes.push_back(bufferType);
+ bufferOffsets.push_back(bufferOffset);
+ bufferLengths.push_back(bufferLength);
+ } while (succeeded(parser.parseOptionalComma()));
+ return success();
+}
+
+static void printDispatchBindings(OpAsmPrinter &p, Operation *op,
+ ValueRange buffers, TypeRange bufferTypes,
+ ValueRange bufferOffsets,
+ ValueRange bufferLengths) {
+ llvm::interleaveComma(
+ llvm::zip(buffers, bufferTypes, bufferOffsets, bufferLengths), p,
+ [&](std::tuple<Value, Type, Value, Value> it) {
+ p.printNewline();
+ p << " ";
+ p << "(";
+ p.printOperand(std::get<0>(it));
+ p << " : ";
+ p.printType(std::get<1>(it));
+ p << ")[";
+ p.printOperand(std::get<2>(it));
+ p << ", ";
+ p.printOperand(std::get<3>(it));
+ p << "]";
+ });
+ p.printNewline();
+}
+
+//===----------------------------------------------------------------------===//
+// hal_loader.executable.query_support
+//===----------------------------------------------------------------------===//
+
+void ExecutableQuerySupportOp::getAsmResultNames(
+ function_ref<void(Value, StringRef)> setNameFn) {
+ setNameFn(getSupported(), (getExecutableFormat() + "_supported").str());
+}
+
+//===----------------------------------------------------------------------===//
+// hal_loader.executable.load
+//===----------------------------------------------------------------------===//
+
+void ExecutableLoadOp::getAsmResultNames(
+ function_ref<void(Value, StringRef)> setNameFn) {
+ setNameFn(getResult(), "exe");
+}
+
+//===----------------------------------------------------------------------===//
+// hal_loader.executable.lookup
+//===----------------------------------------------------------------------===//
+
+void ExecutableLookupOp::getAsmResultNames(
+ function_ref<void(Value, StringRef)> setNameFn) {
+ setNameFn(getResult(), "exe");
+}
+
+LogicalResult ExecutableLookupOp::verifySymbolUses(
+ SymbolTableCollection &symbolTable) {
+ Operation *op = getOperation();
+ auto exportOp = symbolTable.lookupNearestSymbolFrom<IREE::HAL::ExecutableOp>(
+ op, getExecutableAttr());
+ if (!exportOp) {
+ return op->emitOpError() << "undefined executable: " << getExecutable();
+ }
+ return success();
+}
+
+//===----------------------------------------------------------------------===//
+// hal_loader.executable.dispatch
+//===----------------------------------------------------------------------===//
+
+static LogicalResult verifyDispatchBindings(Operation *op, OperandRange buffers,
+ OperandRange offsets,
+ OperandRange lengths) {
+ if (buffers.size() != offsets.size() || buffers.size() != lengths.size()) {
+ return op->emitOpError("binding buffers/offsets/lengths must match; have ")
+ << buffers.size() << "/" << offsets.size() << "/" << lengths.size();
+ }
+ return success();
+}
+
+LogicalResult ExecutableDispatchSymbolOp::verifySymbolUses(
+ SymbolTableCollection &symbolTable) {
+ Operation *op = getOperation();
+ auto exportOp =
+ symbolTable.lookupNearestSymbolFrom<IREE::HAL::ExecutableExportOp>(
+ op, getEntryPoint());
+ if (!exportOp) {
+ return op->emitOpError() << "undefined entry point: " << getEntryPoint();
+ }
+ return verifyDispatchBindings(getOperation(), getBindingBuffers(),
+ getBindingOffsets(), getBindingLengths());
+}
+
+LogicalResult ExecutableDispatchOp::verify() {
+ return verifyDispatchBindings(getOperation(), getBindingBuffers(),
+ getBindingOffsets(), getBindingLengths());
+}
+
+namespace {
+
+// Folds subspan ranges into dispatch resource ranges.
+struct FoldBindingSubspansIntoDispatchOp
+ : public OpRewritePattern<ExecutableDispatchOp> {
+ using OpRewritePattern::OpRewritePattern;
+ LogicalResult matchAndRewrite(ExecutableDispatchOp op,
+ PatternRewriter &rewriter) const override {
+ bool didChangeAny = false;
+ SmallVector<Value> bindingBuffers;
+ SmallVector<Value> bindingOffsets;
+ SmallVector<Value> bindingLengths;
+ for (auto [bindingBuffer, bindingOffset] :
+ llvm::zip(op.getBindingBuffers(), op.getBindingOffsets())) {
+ auto subspanOp =
+ IREE::Util::BufferSubspanOp::findSubspanOp(bindingBuffer);
+ if (!subspanOp) {
+ // No subspan, unchanged.
+ bindingBuffers.push_back(bindingBuffer);
+ bindingOffsets.push_back(bindingOffset);
+ continue;
+ }
+ // Update storage to the source of the subspan and add the subspan offset.
+ didChangeAny = true;
+ auto fusedLoc = rewriter.getFusedLoc({subspanOp.getLoc(), op.getLoc()});
+ auto newOffset = rewriter.createOrFold<arith::AddIOp>(
+ fusedLoc, subspanOp.getSourceOffset(), bindingOffset);
+ bindingBuffers.push_back(subspanOp.getSource());
+ bindingOffsets.push_back(newOffset);
+ }
+ if (!didChangeAny) return failure();
+ rewriter.updateRootInPlace(op, [&]() {
+ op.getBindingBuffersMutable().assign(bindingBuffers);
+ op.getBindingOffsetsMutable().assign(bindingOffsets);
+ });
+ return success();
+ }
+};
+
+} // namespace
+
+void ExecutableDispatchOp::getCanonicalizationPatterns(
+ RewritePatternSet &results, MLIRContext *context) {
+ results.insert<FoldBindingSubspansIntoDispatchOp>(context);
+}
+
+} // namespace Loader
+} // namespace HAL
+} // namespace IREE
+} // namespace iree_compiler
+} // namespace mlir
+
+//===----------------------------------------------------------------------===//
+// TableGen definitions (intentionally last)
+//===----------------------------------------------------------------------===//
+
+#define GET_OP_CLASSES
+#include "iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderOps.cpp.inc"
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderOps.h b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderOps.h
new file mode 100644
index 0000000..7c5776b
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderOps.h
@@ -0,0 +1,26 @@
+// Copyright 2022 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
+
+#ifndef IREE_COMPILER_DIALECT_MODULES_HAL_LOADER_IR_HALLOADEROPS_H_
+#define IREE_COMPILER_DIALECT_MODULES_HAL_LOADER_IR_HALLOADEROPS_H_
+
+#include <cstdint>
+
+#include "iree/compiler/Dialect/HAL/IR/HALTypes.h"
+#include "iree/compiler/Dialect/Util/IR/UtilTraits.h"
+#include "mlir/IR/Attributes.h"
+#include "mlir/IR/BuiltinOps.h"
+#include "mlir/IR/BuiltinTypes.h"
+#include "mlir/IR/Dialect.h"
+#include "mlir/IR/OpDefinition.h"
+#include "mlir/IR/OpImplementation.h"
+#include "mlir/IR/SymbolTable.h"
+#include "mlir/Interfaces/SideEffectInterfaces.h"
+
+#define GET_OP_CLASSES
+#include "iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderOps.h.inc" // IWYU pragma: keep
+
+#endif // IREE_COMPILER_DIALECT_MODULES_HAL_LOADER_IR_HALLOADEROPS_H_
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderOps.td b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderOps.td
new file mode 100644
index 0000000..24866e3
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderOps.td
@@ -0,0 +1,188 @@
+// Copyright 2022 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
+
+#ifndef IREE_DIALECT_MODULES_HAL_LOADER_OPS
+#define IREE_DIALECT_MODULES_HAL_LOADER_OPS
+
+include "iree/compiler/Dialect/HAL/IR/HALBase.td"
+include "iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderBase.td"
+include "iree/compiler/Dialect/Util/IR/UtilAttrs.td"
+include "iree/compiler/Dialect/Util/IR/UtilInterfaces.td"
+include "mlir/IR/OpAsmInterface.td"
+include "mlir/IR/SymbolInterfaces.td"
+include "mlir/Interfaces/SideEffectInterfaces.td"
+
+class HALLoader_PureOp<string mnemonic, list<Trait> traits = []> :
+ HALLoader_Op<mnemonic, !listconcat(traits, [NoSideEffect])>;
+
+//===----------------------------------------------------------------------===//
+// !hal.executable / iree_hal_executable_t
+//===----------------------------------------------------------------------===//
+
+def HALLoader_ExecutableQuerySupportOp :
+ HALLoader_PureOp<"executable.query_support", [
+ DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>,
+ ]> {
+ let summary = "queries whether an executable format is supported";
+ let description = [{
+ Returns true if the given format is supported by the device loader. This
+ does not guarantee that loading will succeed as the executable may require
+ functionality that cannot be met my the hosting runtime environment.
+ }];
+
+ let arguments = (ins
+ StrAttr:$executable_format
+ );
+ let results = (outs
+ I1:$supported
+ );
+
+ let assemblyFormat = [{
+ `format` `(` $executable_format `)`
+ `:` type($supported)
+ attr-dict-with-keyword
+ }];
+}
+
+def HALLoader_ExecutableLoadOp :
+ HALLoader_PureOp<"executable.load", [
+ DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>,
+ ]> {
+ let summary = "dynamically loads an executable";
+ let description = [{
+ Creates, loads, and dynamically links an executable.
+
+ Optional constants provide for specialization of the executable based on
+ runtime-derived parameters.
+ }];
+
+ let arguments = (ins
+ StrAttr:$format,
+ Util_BufferType:$data,
+ Variadic<I32>:$constants
+ );
+ let results = (outs
+ HAL_Executable:$result
+ );
+
+ let assemblyFormat = [{
+ `format` `(` $format `)`
+ `data` `(` $data `)`
+ (`constants` `(` `[` $constants^ `]` `)`)?
+ `:` type($result)
+ attr-dict-with-keyword
+ }];
+}
+
+def HALLoader_ExecutableLookupOp :
+ HALLoader_PureOp<"executable.lookup", [
+ DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>,
+ DeclareOpInterfaceMethods<SymbolUserOpInterface>,
+ ]> {
+ let summary = [{executable cache lookup pseudo-op}];
+ let description = [{
+ Used during conversion to provide a placeholder for a globally cached and
+ possibly lazy-initialized executable.
+ }];
+
+ let arguments = (ins
+ FlatSymbolRefAttr:$executable
+ );
+ let results = (outs
+ HAL_Executable:$result
+ );
+
+ let assemblyFormat = [{
+ `executable` `(` $executable `)`
+ `:` type($result)
+ attr-dict-with-keyword
+ }];
+}
+
+def HALLoader_ExecutableDispatchSymbolOp :
+ HALLoader_Op<"executable.dispatch.symbol", [
+ AttrSizedOperandSegments,
+ DeclareOpInterfaceMethods<SymbolUserOpInterface>,
+ ]> {
+ let summary = [{inline executable dispatch operation}];
+ let description = [{
+ Dispatches execution to an executable entry point with the given parameters.
+ The entry point is a symbolic reference to an exported entry point.
+ }];
+
+ let arguments = (ins
+ HAL_Executable:$executable,
+ SymbolRefAttr:$entry_point,
+ HAL_Dim:$workgroup_x,
+ HAL_Dim:$workgroup_y,
+ HAL_Dim:$workgroup_z,
+ Variadic<I32>:$push_constants,
+ Variadic<Util_BufferType>:$binding_buffers,
+ Variadic<HAL_DeviceSize>:$binding_offsets,
+ Variadic<HAL_DeviceSize>:$binding_lengths
+ );
+
+ let assemblyFormat = [{
+ `executable` `(` $executable `:` type($executable) `)`
+ `target` `(` $entry_point `)`
+ `workgroups` `(` `[`
+ $workgroup_x `,`
+ $workgroup_y `,`
+ $workgroup_z
+ `]` `)`
+ (`constants` `(` `[` $push_constants^ `]` `)`)?
+ `bindings` `(` `[`
+ custom<DispatchBindings>($binding_buffers,
+ type($binding_buffers),
+ $binding_offsets,
+ $binding_lengths)
+ `]` `)`
+ attr-dict-with-keyword
+ }];
+}
+
+def HALLoader_ExecutableDispatchOp :
+ HALLoader_Op<"executable.dispatch", [AttrSizedOperandSegments]> {
+ let summary = [{inline executable dispatch operation}];
+ let description = [{
+ Dispatches execution to an executable entry point with the given parameters.
+ }];
+
+ let arguments = (ins
+ HAL_Executable:$executable,
+ HAL_OrdinalAttr:$entry_point,
+ HAL_Dim:$workgroup_x,
+ HAL_Dim:$workgroup_y,
+ HAL_Dim:$workgroup_z,
+ Variadic<I32>:$push_constants,
+ Variadic<Util_BufferType>:$binding_buffers,
+ Variadic<HAL_DeviceSize>:$binding_offsets,
+ Variadic<HAL_DeviceSize>:$binding_lengths
+ );
+
+ let assemblyFormat = [{
+ `executable` `(` $executable `:` type($executable) `)`
+ `` `[` $entry_point `]`
+ `workgroups` `(` `[`
+ $workgroup_x `,`
+ $workgroup_y `,`
+ $workgroup_z
+ `]` `)`
+ (`constants` `(` `[` $push_constants^ `]` `)`)?
+ `bindings` `(` `[`
+ custom<DispatchBindings>($binding_buffers,
+ type($binding_buffers),
+ $binding_offsets,
+ $binding_lengths)
+ `]` `)`
+ attr-dict-with-keyword
+ }];
+
+ let hasVerifier = 1;
+ let hasCanonicalizer = 1;
+}
+
+#endif // IREE_DIALECT_MODULES_HAL_LOADER_OPS
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/test/BUILD b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/test/BUILD
new file mode 100644
index 0000000..b3b747d
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/test/BUILD
@@ -0,0 +1,28 @@
+# Copyright 2022 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
+
+load("//build_tools/bazel:iree_lit_test.bzl", "iree_lit_test_suite")
+load("//build_tools/bazel:enforce_glob.bzl", "enforce_glob")
+
+package(
+ features = ["layering_check"],
+ licenses = ["notice"], # Apache 2.0
+)
+
+iree_lit_test_suite(
+ name = "lit",
+ srcs = enforce_glob(
+ [
+ "dispatch_folding.mlir",
+ ],
+ include = ["*.mlir"],
+ ),
+ cfg = "//compiler:lit.cfg.py",
+ tools = [
+ "//tools:iree-opt",
+ "@llvm-project//llvm:FileCheck",
+ ],
+)
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/test/CMakeLists.txt b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/test/CMakeLists.txt
new file mode 100644
index 0000000..7779a10
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/test/CMakeLists.txt
@@ -0,0 +1,23 @@
+################################################################################
+# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from #
+# compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/test/BUILD #
+# #
+# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary #
+# CMake-only content. #
+# #
+# To disable autogeneration for this file entirely, delete this header. #
+################################################################################
+
+iree_add_all_subdirs()
+
+iree_lit_test_suite(
+ NAME
+ lit
+ SRCS
+ "dispatch_folding.mlir"
+ TOOLS
+ FileCheck
+ iree-opt
+)
+
+### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ###
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/test/dispatch_folding.mlir b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/test/dispatch_folding.mlir
new file mode 100644
index 0000000..af29692
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR/test/dispatch_folding.mlir
@@ -0,0 +1,30 @@
+// RUN: iree-opt --split-input-file --canonicalize -cse %s | iree-opt --allow-unregistered-dialect --split-input-file | FileCheck %s
+
+// CHECK-LABEL: @fold_binding_subspans_into_dispatch
+func.func @fold_binding_subspans_into_dispatch(
+ // CHECK-SAME: %[[EXECUTABLE:.+]]: !hal.executable,
+ %executable: !hal.executable,
+ // CHECK-SAME: %[[BUFFER:.+]]: !util.buffer, %[[SUBSPAN_OFFSET:.+]]: index, %[[SUBSPAN_LENGTH:.+]]: index
+ %buffer: !util.buffer, %subspan_offset: index, %subspan_length: index) {
+ %c1 = arith.constant 1 : index
+
+ %buffer_size = util.buffer.size %buffer : !util.buffer
+ %subspan = util.buffer.subspan %buffer[%subspan_offset] : !util.buffer{%buffer_size} -> !util.buffer{%subspan_length}
+
+ // CHECK-DAG: %[[BINDING_OFFSET:.+]] = arith.constant 100
+ %binding_offset = arith.constant 100 : index
+ // CHECK-DAG: %[[BINDING_LENGTH:.+]] = arith.constant 128
+ %binding_length = arith.constant 128 : index
+
+ // CHECK-DAG: %[[ABSOLUTE_OFFSET:.+]] = arith.addi %[[SUBSPAN_OFFSET]], %[[BINDING_OFFSET]] : index
+
+ // CHECK: hal_loader.executable.dispatch
+ hal_loader.executable.dispatch
+ executable(%executable : !hal.executable)[16]
+ workgroups([%c1, %c1, %c1])
+ bindings([
+ // CHECK: (%[[BUFFER]] : !util.buffer)[%[[ABSOLUTE_OFFSET]], %[[BINDING_LENGTH]]]
+ (%subspan : !util.buffer)[%binding_offset, %binding_length]
+ ])
+ return
+}
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/BUILD b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/BUILD
new file mode 100644
index 0000000..da94dd0
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/BUILD
@@ -0,0 +1,88 @@
+# Copyright 2022 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
+
+load("//build_tools/bazel:build_defs.oss.bzl", "iree_compiler_cc_library")
+load("//build_tools/bazel:iree_tablegen.bzl", "iree_gentbl_cc_library")
+
+package(
+ default_visibility = ["//visibility:public"],
+ features = ["layering_check"],
+ licenses = ["notice"], # Apache 2.0
+)
+
+iree_compiler_cc_library(
+ name = "Transforms",
+ srcs = [
+ "Conversion.cpp",
+ "MaterializeExecutables.cpp",
+ "Passes.cpp",
+ "ResolveExportOrdinals.cpp",
+ ],
+ hdrs = ["Passes.h"],
+ deps = [
+ ":PassHeaders",
+ "//compiler/src/iree/compiler/Dialect/HAL/Conversion/StandardToHAL",
+ "//compiler/src/iree/compiler/Dialect/HAL/Conversion/UtilToHAL",
+ "//compiler/src/iree/compiler/Dialect/HAL/IR:HALDialect",
+ "//compiler/src/iree/compiler/Dialect/HAL/Target",
+ "//compiler/src/iree/compiler/Dialect/HAL/Transforms",
+ "//compiler/src/iree/compiler/Dialect/Modules/HAL/Inline/Conversion/HALToHALInline",
+ "//compiler/src/iree/compiler/Dialect/Modules/HAL/Inline/Conversion/StreamToHALInline",
+ "//compiler/src/iree/compiler/Dialect/Modules/HAL/Inline/IR",
+ "//compiler/src/iree/compiler/Dialect/Modules/HAL/Inline/IR:HALInlineDialect",
+ "//compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Conversion/StreamToHALLoader",
+ "//compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR",
+ "//compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR:HALLoaderDialect",
+ "//compiler/src/iree/compiler/Dialect/Stream/IR",
+ "//compiler/src/iree/compiler/Dialect/Util/Conversion",
+ "//compiler/src/iree/compiler/Dialect/Util/IR",
+ "//compiler/src/iree/compiler/Dialect/Util/Transforms",
+ "//compiler/src/iree/compiler/Utils",
+ "@llvm-project//llvm:Support",
+ "@llvm-project//mlir:AffineDialect",
+ "@llvm-project//mlir:ArithmeticDialect",
+ "@llvm-project//mlir:ArithmeticTransforms",
+ "@llvm-project//mlir:ControlFlowDialect",
+ "@llvm-project//mlir:FuncDialect",
+ "@llvm-project//mlir:FuncTransforms",
+ "@llvm-project//mlir:IR",
+ "@llvm-project//mlir:MathDialect",
+ "@llvm-project//mlir:MathTransforms",
+ "@llvm-project//mlir:MemRefDialect",
+ "@llvm-project//mlir:Pass",
+ "@llvm-project//mlir:SCFDialect",
+ "@llvm-project//mlir:SCFToControlFlow",
+ "@llvm-project//mlir:Support",
+ "@llvm-project//mlir:Transforms",
+ ],
+)
+
+iree_compiler_cc_library(
+ name = "PassHeaders",
+ hdrs = [
+ "PassDetail.h",
+ "Passes.h",
+ "Passes.h.inc",
+ ],
+ deps = [
+ ":PassesIncGen",
+ "@llvm-project//mlir:Pass",
+ "@llvm-project//mlir:Transforms",
+ ],
+)
+
+iree_gentbl_cc_library(
+ name = "PassesIncGen",
+ tbl_outs = [
+ (
+ ["--gen-pass-decls"],
+ "Passes.h.inc",
+ ),
+ ],
+ tblgen = "@llvm-project//mlir:mlir-tblgen",
+ td_file = "Passes.td",
+ deps = ["@llvm-project//mlir:PassBaseTdFiles"],
+)
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/CMakeLists.txt b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/CMakeLists.txt
new file mode 100644
index 0000000..245bcec
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/CMakeLists.txt
@@ -0,0 +1,84 @@
+################################################################################
+# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from #
+# compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/BUILD #
+# #
+# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary #
+# CMake-only content. #
+# #
+# To disable autogeneration for this file entirely, delete this header. #
+################################################################################
+
+iree_add_all_subdirs()
+
+iree_cc_library(
+ NAME
+ Transforms
+ HDRS
+ "Passes.h"
+ SRCS
+ "Conversion.cpp"
+ "MaterializeExecutables.cpp"
+ "Passes.cpp"
+ "ResolveExportOrdinals.cpp"
+ DEPS
+ ::PassHeaders
+ LLVMSupport
+ MLIRAffineDialect
+ MLIRArithmeticDialect
+ MLIRArithmeticTransforms
+ MLIRControlFlowDialect
+ MLIRFuncDialect
+ MLIRFuncTransforms
+ MLIRIR
+ MLIRMathDialect
+ MLIRMathTransforms
+ MLIRMemRefDialect
+ MLIRPass
+ MLIRSCFDialect
+ MLIRSCFToControlFlow
+ MLIRSupport
+ MLIRTransforms
+ iree::compiler::Dialect::HAL::Conversion::StandardToHAL
+ iree::compiler::Dialect::HAL::Conversion::UtilToHAL
+ iree::compiler::Dialect::HAL::IR::HALDialect
+ iree::compiler::Dialect::HAL::Target
+ iree::compiler::Dialect::HAL::Transforms
+ iree::compiler::Dialect::Modules::HAL::Inline::Conversion::HALToHALInline
+ iree::compiler::Dialect::Modules::HAL::Inline::Conversion::StreamToHALInline
+ iree::compiler::Dialect::Modules::HAL::Inline::IR
+ iree::compiler::Dialect::Modules::HAL::Inline::IR::HALInlineDialect
+ iree::compiler::Dialect::Modules::HAL::Loader::Conversion::StreamToHALLoader
+ iree::compiler::Dialect::Modules::HAL::Loader::IR
+ iree::compiler::Dialect::Modules::HAL::Loader::IR::HALLoaderDialect
+ iree::compiler::Dialect::Stream::IR
+ iree::compiler::Dialect::Util::Conversion
+ iree::compiler::Dialect::Util::IR
+ iree::compiler::Dialect::Util::Transforms
+ iree::compiler::Utils
+ PUBLIC
+)
+
+iree_cc_library(
+ NAME
+ PassHeaders
+ HDRS
+ "PassDetail.h"
+ "Passes.h"
+ "Passes.h.inc"
+ DEPS
+ ::PassesIncGen
+ MLIRPass
+ MLIRTransforms
+ PUBLIC
+)
+
+iree_tablegen_library(
+ NAME
+ PassesIncGen
+ TD_FILE
+ "Passes.td"
+ OUTS
+ --gen-pass-decls Passes.h.inc
+)
+
+### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ###
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/Conversion.cpp b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/Conversion.cpp
new file mode 100644
index 0000000..b94a529
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/Conversion.cpp
@@ -0,0 +1,111 @@
+// Copyright 2022 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 "iree/compiler/Dialect/HAL/Conversion/StandardToHAL/ConvertStandardToHAL.h"
+#include "iree/compiler/Dialect/HAL/Conversion/UtilToHAL/ConvertUtilToHAL.h"
+#include "iree/compiler/Dialect/HAL/IR/HALDialect.h"
+#include "iree/compiler/Dialect/Modules/HAL/Inline/Conversion/HALToHALInline/ConvertHALToHALInline.h"
+#include "iree/compiler/Dialect/Modules/HAL/Inline/Conversion/StreamToHALInline/ConvertStreamToHALInline.h"
+#include "iree/compiler/Dialect/Modules/HAL/Inline/IR/HALInlineDialect.h"
+#include "iree/compiler/Dialect/Modules/HAL/Loader/Conversion/StreamToHALLoader/ConvertStreamToHALLoader.h"
+#include "iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderDialect.h"
+#include "iree/compiler/Dialect/Modules/HAL/Loader/Transforms/PassDetail.h"
+#include "iree/compiler/Dialect/Modules/HAL/Loader/Transforms/Passes.h"
+#include "iree/compiler/Dialect/Stream/IR/StreamDialect.h"
+#include "iree/compiler/Dialect/Util/Conversion/ConversionPatterns.h"
+#include "iree/compiler/Dialect/Util/IR/UtilDialect.h"
+#include "llvm/ADT/STLExtras.h"
+#include "mlir/Dialect/Affine/IR/AffineOps.h"
+#include "mlir/Dialect/Func/IR/FuncOps.h"
+#include "mlir/Dialect/Math/IR/Math.h"
+#include "mlir/Dialect/SCF/IR/SCF.h"
+#include "mlir/IR/BuiltinOps.h"
+#include "mlir/IR/PatternMatch.h"
+#include "mlir/Pass/Pass.h"
+#include "mlir/Pass/PassRegistry.h"
+#include "mlir/Transforms/DialectConversion.h"
+
+namespace mlir {
+namespace iree_compiler {
+namespace IREE {
+namespace HAL {
+namespace Loader {
+
+// Runs conversion with registered input dialects.
+class ConversionPass : public ConversionBase<ConversionPass> {
+ public:
+ void getDependentDialects(DialectRegistry ®istry) const override {
+ registry.insert<IREE::Util::UtilDialect, IREE::HAL::HALDialect,
+ IREE::HAL::Inline::HALInlineDialect,
+ IREE::HAL::Loader::HALLoaderDialect,
+ mlir::arith::ArithmeticDialect>();
+ }
+
+ void runOnOperation() override {
+ auto *context = &getContext();
+
+ // Ensure all input dialects go away.
+ ConversionTarget conversionTarget(*context);
+ conversionTarget
+ .addLegalDialect<mlir::func::FuncDialect, mlir::scf::SCFDialect,
+ mlir::arith::ArithmeticDialect>();
+
+ TypeConverter typeConverter;
+ RewritePatternSet patterns(context);
+
+ // Pass-through.
+ typeConverter.addConversion([](IndexType type) { return type; });
+ typeConverter.addConversion([](IntegerType type) { return type; });
+ typeConverter.addConversion([](FloatType type) { return type; });
+ typeConverter.addConversion(
+ [](IREE::Util::BufferType type) { return type; });
+
+ // Convert stream into `hal_inline`, which mostly entails ignoring ops.
+ // We override those related to executables to `hal_loader` by way of high
+ // pattern benefits.
+ conversionTarget.addLegalDialect<IREE::HAL::Inline::HALInlineDialect>();
+ populateStreamToHALInlinePatterns(context, conversionTarget, typeConverter,
+ patterns);
+ conversionTarget.addLegalDialect<IREE::HAL::Loader::HALLoaderDialect>();
+ populateStreamToHALLoaderPatterns(context, conversionTarget, typeConverter,
+ patterns);
+
+ // Convert some common things into HAL, reusing those conversions.
+ populateUtilToHALPatterns(context, conversionTarget, typeConverter,
+ patterns);
+ populateStandardToHALPatterns(context, conversionTarget, typeConverter,
+ patterns);
+
+ // Convert any full `hal` ops into `hal_inline` ops.
+ conversionTarget.addIllegalDialect<IREE::HAL::HALDialect>();
+ populateHALToHALInlinePatterns(context, conversionTarget, typeConverter,
+ patterns);
+
+ // Generic conversion.
+ conversionTarget.addLegalDialect<IREE::Util::UtilDialect>();
+ populateUtilConversionPatterns(context, conversionTarget, typeConverter,
+ patterns);
+ populateGenericStructuralConversionPatterns(context, conversionTarget,
+ typeConverter, patterns);
+
+ if (failed(applyPartialConversion(getOperation(), conversionTarget,
+ std::move(patterns)))) {
+ getOperation().emitError()
+ << "conversion to the hal_inline + hal_loader dialects failed";
+ return signalPassFailure();
+ }
+ }
+};
+
+std::unique_ptr<OperationPass<mlir::ModuleOp>> createConversionPass() {
+ return std::make_unique<ConversionPass>();
+}
+
+} // namespace Loader
+} // namespace HAL
+} // namespace IREE
+} // namespace iree_compiler
+} // namespace mlir
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/MaterializeExecutables.cpp b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/MaterializeExecutables.cpp
new file mode 100644
index 0000000..c62ea78
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/MaterializeExecutables.cpp
@@ -0,0 +1,173 @@
+// Copyright 2022 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 "iree/compiler/Dialect/HAL/IR/HALDialect.h"
+#include "iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderDialect.h"
+#include "iree/compiler/Dialect/Modules/HAL/Loader/Transforms/PassDetail.h"
+#include "iree/compiler/Dialect/Modules/HAL/Loader/Transforms/Passes.h"
+#include "iree/compiler/Dialect/Util/IR/UtilDialect.h"
+#include "mlir/Dialect/Arithmetic/IR/Arithmetic.h"
+#include "mlir/Dialect/ControlFlow/IR/ControlFlow.h"
+#include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h"
+#include "mlir/Dialect/Func/IR/FuncOps.h"
+#include "mlir/IR/BuiltinOps.h"
+#include "mlir/Pass/Pass.h"
+#include "mlir/Pass/PassRegistry.h"
+
+namespace mlir {
+namespace iree_compiler {
+namespace IREE {
+namespace HAL {
+namespace Loader {
+
+static void replaceExecutableWithGlobal(IREE::HAL::ExecutableOp executableOp) {
+ OpBuilder moduleBuilder(executableOp);
+
+ auto loc = executableOp.getLoc();
+
+ // Create global representing the loaded executable.
+ // This matches the executable name and is used to directly access the
+ // executable reference during dispatches.
+ auto executableType = moduleBuilder.getType<IREE::HAL::ExecutableType>();
+ auto globalOp = moduleBuilder.create<IREE::Util::GlobalOp>(
+ loc, executableOp.getName(), /*isMutable=*/false, executableType);
+ globalOp.setPrivate();
+
+ // Create initializer that selects the right binary and loads it.
+ auto initializerOp = moduleBuilder.create<IREE::Util::InitializerOp>(loc);
+ auto entryBuilder = OpBuilder::atBlockBegin(initializerOp.addEntryBlock());
+
+ // Reserve one block per attempt to load a binary.
+ auto binaryOps =
+ llvm::to_vector(executableOp.getOps<IREE::HAL::ExecutableBinaryOp>());
+ SmallVector<Block *> queryBlocks;
+ SmallVector<Block *> loadBlocks;
+ for (size_t i = 0; i < binaryOps.size(); ++i) {
+ queryBlocks.push_back(initializerOp.addBlock());
+ loadBlocks.push_back(initializerOp.addBlock());
+ }
+
+ // Failure block when no binary is supported.
+ auto *failBlock = initializerOp.addBlock();
+ {
+ auto failBuilder = OpBuilder::atBlockBegin(failBlock);
+ Value status = failBuilder.create<arith::ConstantIntOp>(
+ loc, static_cast<int>(IREE::Util::StatusCode::Unavailable), 32);
+ failBuilder.create<IREE::Util::StatusCheckOkOp>(
+ loc, status,
+ "none of the executable binaries in the module are supported by the "
+ "runtime");
+ failBuilder.create<IREE::Util::InitializerReturnOp>(loc);
+ }
+
+ // Exit block takes the loaded executable and stores it.
+ auto *exitBlock = initializerOp.addBlock();
+ {
+ auto exitBuilder = OpBuilder::atBlockBegin(exitBlock);
+ auto executableArg = exitBlock->addArgument(executableType, loc);
+ exitBuilder.create<IREE::Util::GlobalStoreOp>(loc, executableArg,
+ globalOp.getName());
+ exitBuilder.create<IREE::Util::InitializerReturnOp>(loc);
+ }
+
+ // Start with the first try.
+ if (!queryBlocks.empty()) {
+ entryBuilder.create<cf::BranchOp>(loc, queryBlocks[0]);
+ } else {
+ entryBuilder.create<cf::BranchOp>(loc, failBlock);
+ }
+
+ // Build the full chain of try ops. An scf.switch would be nice...
+ // We could also avoid this by having an op that given a list of formats
+ // selected the ones that were supported - that'd result in smaller binary
+ // sizes but not allow for customization of selection logic. Today this
+ // looks bad because our selection logic is dumb :)
+ //
+ // ^queryBlock:
+ // %supported = executable.query_support "format" : i1
+ // cond_br %supported, ^loadBlock, ^nextBlock
+ // ^loadBlock:
+ // %exe = executable.load : !hal.executable
+ // br ^exit(%exe)
+ // ^nextBlock: ...
+ for (unsigned i = 0; i < binaryOps.size(); ++i) {
+ auto binaryOp = binaryOps[i];
+ auto binaryLoc = binaryOp.getLoc();
+
+ // Query whether the format is supported and branch to the load block if
+ // it is. Otherwise we go to the next query block or fail if at the end.
+ auto queryBuilder = OpBuilder::atBlockBegin(queryBlocks[i]);
+ auto *nextBlock = i + 1 < binaryOps.size() ? queryBlocks[i + 1] : failBlock;
+ Value isSupported =
+ queryBuilder.create<IREE::HAL::Loader::ExecutableQuerySupportOp>(
+ binaryLoc, queryBuilder.getI1Type(), binaryOp.getFormatAttr());
+ queryBuilder.create<cf::CondBranchOp>(binaryLoc, isSupported, loadBlocks[i],
+ ValueRange{}, nextBlock,
+ ValueRange{});
+
+ // Load the executable. This may still fail but it'll propagate the error
+ // up to the user with the full status message instead of continuing
+ // execution.
+ auto loadBuilder = OpBuilder::atBlockBegin(loadBlocks[i]);
+ auto alignmentAttr = loadBuilder.getIndexAttr(64);
+ Value binaryData = loadBuilder.create<IREE::Util::BufferConstantOp>(
+ binaryLoc, binaryOp.getNameAttr(), binaryOp.getData(), alignmentAttr,
+ binaryOp.getMimeTypeAttr());
+ SmallVector<Value> constants; // TBD
+ Value executable = loadBuilder.create<IREE::HAL::Loader::ExecutableLoadOp>(
+ binaryLoc, executableType, binaryOp.getFormatAttr(), binaryData,
+ constants);
+ loadBuilder.create<cf::BranchOp>(binaryLoc, exitBlock,
+ ValueRange{executable});
+ }
+
+ // Op goes away to get replaced with a global.
+ executableOp.erase();
+}
+
+// Runs conversion with registered input dialects.
+class MaterializeExecutablesPass
+ : public MaterializeExecutablesBase<MaterializeExecutablesPass> {
+ public:
+ void getDependentDialects(DialectRegistry ®istry) const override {
+ registry.insert<IREE::Util::UtilDialect, IREE::HAL::HALDialect,
+ IREE::HAL::Loader::HALLoaderDialect,
+ arith::ArithmeticDialect, cf::ControlFlowDialect>();
+ }
+
+ void runOnOperation() override {
+ mlir::ModuleOp moduleOp = getOperation();
+
+ // Walk executables and convert each one to a global.
+ for (auto executableOp : llvm::make_early_inc_range(
+ moduleOp.getOps<IREE::HAL::ExecutableOp>())) {
+ replaceExecutableWithGlobal(executableOp);
+ }
+
+ // Find lookup ops referencing an executable and swap it to a global load.
+ for (auto funcOp : llvm::make_early_inc_range(
+ moduleOp.getOps<mlir::FunctionOpInterface>())) {
+ funcOp.walk([&](IREE::HAL::Loader::ExecutableLookupOp lookupOp) {
+ Value executable = OpBuilder(lookupOp).create<IREE::Util::GlobalLoadOp>(
+ lookupOp.getLoc(), lookupOp.getResult().getType(),
+ lookupOp.getExecutableAttr());
+ lookupOp.replaceAllUsesWith(executable);
+ lookupOp.erase();
+ });
+ }
+ }
+};
+
+std::unique_ptr<OperationPass<mlir::ModuleOp>>
+createMaterializeExecutablesPass() {
+ return std::make_unique<MaterializeExecutablesPass>();
+}
+
+} // namespace Loader
+} // namespace HAL
+} // namespace IREE
+} // namespace iree_compiler
+} // namespace mlir
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/PassDetail.h b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/PassDetail.h
new file mode 100644
index 0000000..9ecc064
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/PassDetail.h
@@ -0,0 +1,29 @@
+// Copyright 2022 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
+
+#ifndef IREE_COMPILER_DIALECT_MODULES_HAL_LOADER_TRANSFORMS_PASS_DETAIL_H_
+#define IREE_COMPILER_DIALECT_MODULES_HAL_LOADER_TRANSFORMS_PASS_DETAIL_H_
+
+#include "iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderOps.h"
+#include "mlir/Dialect/Func/IR/FuncOps.h"
+#include "mlir/Pass/Pass.h"
+
+namespace mlir {
+namespace iree_compiler {
+namespace IREE {
+namespace HAL {
+namespace Loader {
+
+#define GEN_PASS_CLASSES
+#include "iree/compiler/Dialect/Modules/HAL/Loader/Transforms/Passes.h.inc"
+
+} // namespace Loader
+} // namespace HAL
+} // namespace IREE
+} // namespace iree_compiler
+} // namespace mlir
+
+#endif // IREE_COMPILER_DIALECT_MODULES_HAL_LOADER_TRANSFORMS_PASS_DETAIL_H_
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/Passes.cpp b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/Passes.cpp
new file mode 100644
index 0000000..afd4763
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/Passes.cpp
@@ -0,0 +1,135 @@
+// Copyright 2022 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 "iree/compiler/Dialect/Modules/HAL/Loader/Transforms/Passes.h"
+
+#include <memory>
+
+#include "iree/compiler/Dialect/HAL/Transforms/Passes.h"
+#include "iree/compiler/Dialect/Util/Transforms/Passes.h"
+#include "iree/compiler/Utils/PassUtils.h"
+#include "mlir/Dialect/Arithmetic/Transforms/Passes.h"
+#include "mlir/Dialect/Func/Transforms/Passes.h"
+#include "mlir/Pass/PassRegistry.h"
+#include "mlir/Transforms/Passes.h"
+
+namespace mlir {
+namespace iree_compiler {
+namespace IREE {
+namespace HAL {
+namespace Loader {
+
+using FunctionLikeNest = MultiOpNest<func::FuncOp, IREE::Util::InitializerOp>;
+
+//===----------------------------------------------------------------------===//
+// Utilities
+//===----------------------------------------------------------------------===//
+
+static void addCleanupPatterns(OpPassManager &passManager) {
+ // Standard MLIR cleanup.
+ passManager.addPass(mlir::createCanonicalizerPass());
+ passManager.addPass(mlir::createCSEPass());
+
+ FunctionLikeNest(passManager)
+ // Simplify util.global accesses; this can help with data flow tracking as
+ // redundant store-loads are removed.
+ .addPass(IREE::Util::createSimplifyGlobalAccessesPass);
+
+ // Cleanup and canonicalization of util.global (and other util ops).
+ passManager.addPass(IREE::Util::createApplyPatternsPass());
+ passManager.addPass(IREE::Util::createFoldGlobalsPass());
+ passManager.addPass(IREE::Util::createFuseGlobalsPass());
+}
+
+//===----------------------------------------------------------------------===//
+// -iree-hal-inline-dynamic-transformation-pipeline
+//===----------------------------------------------------------------------===//
+
+void buildHALInlineDynamicTransformPassPipeline(
+ OpPassManager &passManager, const TargetOptions &targetOptions) {
+ //----------------------------------------------------------------------------
+ // Device assignment and interface materialization
+ //----------------------------------------------------------------------------
+
+ IREE::HAL::buildHALConfigurationPassPipeline(passManager, targetOptions);
+
+ //----------------------------------------------------------------------------
+ // Executable translation
+ //----------------------------------------------------------------------------
+
+ // Translate each executable variant to its target IR form.
+ // It's extremely important this runs parallelized as it's where a large
+ // majority of our compilation time lives (we invoke LLVM and lld and such).
+ //
+ // After this point the executables are opaque blobs and we cannot change
+ // their interfaces.
+ passManager.addNestedPass<IREE::HAL::ExecutableOp>(
+ IREE::HAL::createTranslateExecutablesPass());
+
+ //----------------------------------------------------------------------------
+ // Conversion
+ //----------------------------------------------------------------------------
+
+ // Convert from stream to hal_inline + hal_loader.
+ passManager.addPass(IREE::HAL::Loader::createConversionPass());
+
+ //----------------------------------------------------------------------------
+ // Executable packing and runtime loading
+ //----------------------------------------------------------------------------
+
+ // Link executables together.
+ passManager.addPass(IREE::HAL::createLinkExecutablesPass());
+
+ // Resolve export ordinals from nested symbol references prior to
+ // serialization.
+ passManager.addPass(IREE::HAL::Loader::createResolveExportOrdinalsPass());
+
+ // Serialize executables to their binary forms.
+ passManager.addNestedPass<IREE::HAL::ExecutableOp>(
+ IREE::HAL::createSerializeExecutablesPass(
+ targetOptions.debugLevel, targetOptions.executableIntermediatesPath,
+ targetOptions.executableBinariesPath));
+
+ // NOTE: symbol DCE will destroy executable target contents.
+ passManager.addPass(mlir::createSymbolDCEPass());
+
+ // Materialize executable globals and initializers that load them.
+ passManager.addPass(IREE::HAL::Loader::createMaterializeExecutablesPass());
+
+ //----------------------------------------------------------------------------
+ // Cleanup and canonicalization
+ //----------------------------------------------------------------------------
+
+ addCleanupPatterns(passManager);
+}
+
+//===----------------------------------------------------------------------===//
+// Registration
+//===----------------------------------------------------------------------===//
+
+namespace {
+#define GEN_PASS_REGISTRATION
+#include "iree/compiler/Dialect/Modules/HAL/Loader/Transforms/Passes.h.inc"
+} // namespace
+
+void registerHALLoaderPasses() {
+ // Generated.
+ registerPasses();
+
+ static PassPipelineRegistration<> transformPassPipeline(
+ "iree-hal-inline-dynamic-transformation-pipeline",
+ "Runs the inline HAL executable loader dialect transformation pipeline",
+ [](OpPassManager &passManager) {
+ buildHALInlineDynamicTransformPassPipeline(
+ passManager, TargetOptions::FromFlags::get());
+ });
+}
+
+} // namespace Loader
+} // namespace HAL
+} // namespace IREE
+} // namespace iree_compiler
+} // namespace mlir
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/Passes.h b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/Passes.h
new file mode 100644
index 0000000..4396977
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/Passes.h
@@ -0,0 +1,68 @@
+// Copyright 2022 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
+
+#ifndef IREE_COMPILER_DIALECT_MODULES_HAL_LOADER_TRANSFORMS_PASSES_H_
+#define IREE_COMPILER_DIALECT_MODULES_HAL_LOADER_TRANSFORMS_PASSES_H_
+
+#include "iree/compiler/Dialect/HAL/Target/TargetBackend.h"
+#include "iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderOps.h"
+#include "llvm/ADT/StringMap.h"
+#include "mlir/IR/BuiltinOps.h"
+#include "mlir/Pass/Pass.h"
+#include "mlir/Pass/PassManager.h"
+#include "mlir/Support/LLVM.h"
+
+namespace mlir {
+namespace iree_compiler {
+namespace IREE {
+namespace HAL {
+namespace Loader {
+
+//===----------------------------------------------------------------------===//
+// Helpers
+//===----------------------------------------------------------------------===//
+
+// Adds a set of passes to the given pass manager that run the required
+// HALLoader transforms in the canonical order.
+//
+// Most translation code should prefer to use this instead of manually adding
+// the passes themselves to ensure that expected pass ordering is observed.
+//
+// The expected usage is:
+// <run conversion from TF/HLO/etc -> flow -> stream>
+// buildHALInlineDynamicTransformPassPipeline & run
+// <serialize VM module>
+void buildHALInlineDynamicTransformPassPipeline(
+ OpPassManager &passManager, const TargetOptions &targetOptions);
+
+//===----------------------------------------------------------------------===//
+// Passes
+//===----------------------------------------------------------------------===//
+
+// Converts from the stream dialect into the hal_inline + hal_loader dialects.
+std::unique_ptr<OperationPass<mlir::ModuleOp>> createConversionPass();
+
+// Materializes executable globals and loader code.
+std::unique_ptr<OperationPass<mlir::ModuleOp>>
+createMaterializeExecutablesPass();
+
+// Resolves dispatch operation target export entry point ordinals.
+std::unique_ptr<OperationPass<mlir::ModuleOp>>
+createResolveExportOrdinalsPass();
+
+//===----------------------------------------------------------------------===//
+// Register all Passes
+//===----------------------------------------------------------------------===//
+
+void registerHALLoaderPasses();
+
+} // namespace Loader
+} // namespace HAL
+} // namespace IREE
+} // namespace iree_compiler
+} // namespace mlir
+
+#endif // IREE_COMPILER_DIALECT_MODULES_HAL_LOADER_TRANSFORMS_PASSES_H_
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/Passes.td b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/Passes.td
new file mode 100644
index 0000000..4457660
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/Passes.td
@@ -0,0 +1,27 @@
+// Copyright 2022 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
+
+#ifndef IREE_MODULES_HAL_LOADER_PASSES
+#define IREE_MODULES_HAL_LOADER_PASSES
+
+include "mlir/Pass/PassBase.td"
+
+def Conversion : Pass<"iree-hal-loader-conversion", "mlir::ModuleOp"> {
+ let summary = "Converts from various dialects to the HAL loader dialect";
+ let constructor = "mlir::iree_compiler::IREE::HAL::Loader::createConversionPass()";
+}
+
+def MaterializeExecutables : Pass<"iree-hal-loader-materialize-executables", "mlir::ModuleOp"> {
+ let summary = "Materializes executable globals and loader code";
+ let constructor = "mlir::iree_compiler::IREE::HAL::Loader::createMaterializeExecutablesPass()";
+}
+
+def ResolveExportOrdinals : Pass<"iree-hal-loader-resolve-export-ordinals", "mlir::ModuleOp"> {
+ let summary = "Resolves dispatch operation target export entry point ordinals";
+ let constructor = "mlir::iree_compiler::IREE::HAL::Loader::createResolveExportOrdinalsPass()";
+}
+
+#endif // IREE_MODULES_HAL_LOADER_PASSES
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/ResolveExportOrdinals.cpp b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/ResolveExportOrdinals.cpp
new file mode 100644
index 0000000..3932a71
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/ResolveExportOrdinals.cpp
@@ -0,0 +1,65 @@
+// Copyright 2022 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 "iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderDialect.h"
+#include "iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderOps.h"
+#include "iree/compiler/Dialect/Modules/HAL/Loader/Transforms/PassDetail.h"
+#include "iree/compiler/Dialect/Modules/HAL/Loader/Transforms/Passes.h"
+#include "mlir/Pass/Pass.h"
+#include "mlir/Transforms/GreedyPatternRewriteDriver.h"
+
+namespace mlir {
+namespace iree_compiler {
+namespace IREE {
+namespace HAL {
+namespace Loader {
+
+struct ResolveExecutableDispatchSymbolOp
+ : public OpRewritePattern<IREE::HAL::Loader::ExecutableDispatchSymbolOp> {
+ using OpRewritePattern::OpRewritePattern;
+ LogicalResult matchAndRewrite(
+ IREE::HAL::Loader::ExecutableDispatchSymbolOp op,
+ PatternRewriter &rewriter) const override {
+ auto symbol = SymbolTable::lookupNearestSymbolFrom(op, op.getEntryPoint());
+ assert(symbol && "missing ExecutableEntryPoint symbol");
+ auto exportOp = cast<IREE::HAL::ExecutableExportOp>(symbol);
+ rewriter.replaceOpWithNewOp<IREE::HAL::Loader::ExecutableDispatchOp>(
+ op, op.getExecutable(), exportOp.getOrdinalAttr(), op.getWorkgroupX(),
+ op.getWorkgroupY(), op.getWorkgroupZ(), op.getPushConstants(),
+ op.getBindingBuffers(), op.getBindingOffsets(), op.getBindingLengths());
+ return success();
+ }
+};
+
+class ResolveExportOrdinalsPass
+ : public ResolveExportOrdinalsBase<ResolveExportOrdinalsPass> {
+ public:
+ void getDependentDialects(DialectRegistry ®istry) const override {
+ registry.insert<IREE::HAL::Loader::HALLoaderDialect>();
+ }
+
+ void runOnOperation() override {
+ MLIRContext *context = &getContext();
+ RewritePatternSet patterns(&getContext());
+ patterns.insert<ResolveExecutableDispatchSymbolOp>(context);
+ if (failed(applyPatternsAndFoldGreedily(getOperation(),
+ std::move(patterns)))) {
+ return signalPassFailure();
+ }
+ }
+};
+
+std::unique_ptr<OperationPass<ModuleOp>> createResolveExportOrdinalsPass() {
+ return std::make_unique<ResolveExportOrdinalsPass>();
+}
+
+static PassRegistration<ResolveExportOrdinalsPass> pass;
+
+} // namespace Loader
+} // namespace HAL
+} // namespace IREE
+} // namespace iree_compiler
+} // namespace mlir
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/test/BUILD b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/test/BUILD
new file mode 100644
index 0000000..9ef60cb
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/test/BUILD
@@ -0,0 +1,28 @@
+# Copyright 2022 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
+
+load("//build_tools/bazel:iree_lit_test.bzl", "iree_lit_test_suite")
+load("//build_tools/bazel:enforce_glob.bzl", "enforce_glob")
+
+package(
+ features = ["layering_check"],
+ licenses = ["notice"], # Apache 2.0
+)
+
+iree_lit_test_suite(
+ name = "lit",
+ srcs = enforce_glob(
+ [
+ "materialize_executables.mlir",
+ ],
+ include = ["*.mlir"],
+ ),
+ cfg = "//compiler:lit.cfg.py",
+ tools = [
+ "//tools:iree-opt",
+ "@llvm-project//llvm:FileCheck",
+ ],
+)
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/test/CMakeLists.txt b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/test/CMakeLists.txt
new file mode 100644
index 0000000..11f6cf0
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/test/CMakeLists.txt
@@ -0,0 +1,23 @@
+################################################################################
+# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from #
+# compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/test/BUILD #
+# #
+# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary #
+# CMake-only content. #
+# #
+# To disable autogeneration for this file entirely, delete this header. #
+################################################################################
+
+iree_add_all_subdirs()
+
+iree_lit_test_suite(
+ NAME
+ lit
+ SRCS
+ "materialize_executables.mlir"
+ TOOLS
+ FileCheck
+ iree-opt
+)
+
+### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ###
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/test/materialize_executables.mlir b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/test/materialize_executables.mlir
new file mode 100644
index 0000000..ce4c6e9
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms/test/materialize_executables.mlir
@@ -0,0 +1,48 @@
+// RUN: iree-opt --split-input-file --iree-hal-loader-materialize-executables %s | FileCheck %s
+
+// Tests that executable binaries get moved to initialized globals and lookups
+// get rewritten to point at the globals. Note that we do 2 to ensure we're
+// enumerating all executables and all lookups.
+
+// CHECK-LABEL: util.global private @ex0 : !hal.executable
+// CHECK: util.initializer {
+// CHECK: %[[FORMAT_SUPPORTED:.+]] = hal_loader.executable.query_support format("embedded-elf-x86_64") : i1
+// CHECK: cf.cond_br %[[FORMAT_SUPPORTED]], ^bb2, ^bb3
+// CHECK: ^bb2:
+// CHECK: %[[BINARY_DATA:.+]] = util.buffer.constant "binary" {alignment = 64 : index, mime_type = "application/x-elf"} : !util.buffer = dense<123> : vector<64xi8>
+// CHECK: %[[EXECUTABLE:.+]] = hal_loader.executable.load format("embedded-elf-x86_64") data(%[[BINARY_DATA]]) : !hal.executable
+// CHECK: cf.br ^bb4(%[[EXECUTABLE]] : !hal.executable)
+// CHECK: ^bb3:
+// CHECK: util.status.check_ok
+// CHECK: util.initializer.return
+// CHECK: ^bb4(%[[STORE_VALUE:.+]]: !hal.executable):
+// CHECK: util.global.store %[[STORE_VALUE]], @ex0 : !hal.executable
+// CHECK: util.initializer.return
+// CHECK: }
+hal.executable private @ex0 {
+ hal.executable.binary public @binary attributes {data = dense<123> : vector<64xi8>, format = "embedded-elf-x86_64", mime_type = "application/x-elf"}
+}
+
+// CHECK-LABEL: @get_ex0
+func.func private @get_ex0() -> !hal.executable {
+ // CHECK: %[[EX0:.+]] = util.global.load @ex0 : !hal.executable
+ %ex0 = hal_loader.executable.lookup executable(@ex0) : !hal.executable
+ // CHECK: return %[[EX0]]
+ return %ex0 : !hal.executable
+}
+
+// CHECK: util.global private @ex1 : !hal.executable
+// CHECK: util.initializer
+// CHECK: hal_loader.executable.load format("embedded-elf-aarch64")
+// CHECK: util.global.store {{.+}}, @ex1
+hal.executable private @ex1 {
+ hal.executable.binary public @binary attributes {data = dense<123> : vector<64xi8>, format = "embedded-elf-aarch64", mime_type = "application/x-elf"}
+}
+
+// CHECK-LABEL: @get_ex1
+func.func private @get_ex1() -> !hal.executable {
+ // CHECK: %[[EX1:.+]] = util.global.load @ex1 : !hal.executable
+ %ex1 = hal_loader.executable.lookup executable(@ex1) : !hal.executable
+ // CHECK: return %[[EX1]]
+ return %ex1 : !hal.executable
+}
diff --git a/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/hal_loader.imports.mlir b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/hal_loader.imports.mlir
new file mode 100644
index 0000000..fc9c538
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/hal_loader.imports.mlir
@@ -0,0 +1,40 @@
+// IREE Inline Hardware Abstraction Layer (HAL) loader module imports.
+// This is only used to dynamically load and dispatch executable libraries.
+//
+// This is embedded in the compiler binary and inserted into any module
+// containing inline HAL loader dialect ops (hal_loader.*) that is lowered to
+// the VM dialect.
+vm.module @hal_loader {
+
+//===----------------------------------------------------------------------===//
+// iree_hal_executable_t
+//===----------------------------------------------------------------------===//
+
+// Queries whether the given executable format is supported.
+vm.import @executable.query_support(
+ %executable_format : !vm.buffer
+) -> i32
+attributes {nosideeffects}
+
+// Creates and dynamically links an executable library.
+vm.import @executable.load(
+ %executable_format : !vm.buffer,
+ %executable_data : !vm.buffer,
+ %constants : !vm.buffer
+) -> !vm.ref<!hal.executable>
+attributes {nosideeffects}
+
+// Dispatches a grid with the given densely-packed and 0-aligned push constants
+// and bindings.
+vm.import @executable.dispatch(
+ %executable : !vm.ref<!hal.executable>,
+ %entry_point : i32,
+ %workgroup_x : i32,
+ %workgroup_y : i32,
+ %workgroup_z : i32,
+ %push_constants : i32 ...,
+ // <buffer, offset, length>
+ %bindings : tuple<!vm.buffer, i64, i64>...
+)
+
+} // module
diff --git a/compiler/src/iree/compiler/Pipelines/BUILD b/compiler/src/iree/compiler/Pipelines/BUILD
index 48030f7..a5751f0 100644
--- a/compiler/src/iree/compiler/Pipelines/BUILD
+++ b/compiler/src/iree/compiler/Pipelines/BUILD
@@ -40,6 +40,7 @@
"//compiler/src/iree/compiler/Dialect/HAL/Conversion/HALToVM",
"//compiler/src/iree/compiler/Dialect/HAL/Transforms",
"//compiler/src/iree/compiler/Dialect/Modules/HAL/Inline/Transforms",
+ "//compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms",
"//compiler/src/iree/compiler/Dialect/Stream/Transforms",
"//compiler/src/iree/compiler/Dialect/Util/Transforms",
"//compiler/src/iree/compiler/Dialect/VM/Conversion",
diff --git a/compiler/src/iree/compiler/Pipelines/CMakeLists.txt b/compiler/src/iree/compiler/Pipelines/CMakeLists.txt
index 824566f..3be2507 100644
--- a/compiler/src/iree/compiler/Pipelines/CMakeLists.txt
+++ b/compiler/src/iree/compiler/Pipelines/CMakeLists.txt
@@ -49,6 +49,7 @@
iree::compiler::Dialect::HAL::Conversion::HALToVM
iree::compiler::Dialect::HAL::Transforms
iree::compiler::Dialect::Modules::HAL::Inline::Transforms
+ iree::compiler::Dialect::Modules::HAL::Loader::Transforms
iree::compiler::Dialect::Stream::Transforms
iree::compiler::Dialect::Util::Transforms
iree::compiler::Dialect::VM::Conversion
diff --git a/compiler/src/iree/compiler/Pipelines/Options.cpp b/compiler/src/iree/compiler/Pipelines/Options.cpp
index 1c68eb2..7d2c372 100644
--- a/compiler/src/iree/compiler/Pipelines/Options.cpp
+++ b/compiler/src/iree/compiler/Pipelines/Options.cpp
@@ -100,7 +100,10 @@
"internally and externally."),
clEnumValN(ExecutionModel::InlineStatic, "inline-static",
"Inline host-local in-process execution with executable "
- "code statically linked into the host program.")),
+ "code statically linked into the host program."),
+ clEnumValN(ExecutionModel::InlineDynamic, "inline-dynamic",
+ "Inline host-local in-process execution using dynamic "
+ "executables.")),
llvm::cl::cat(category));
binder.opt<DumpOutputFormat>(
diff --git a/compiler/src/iree/compiler/Pipelines/Options.h b/compiler/src/iree/compiler/Pipelines/Options.h
index 405ee84..56ddf7f 100644
--- a/compiler/src/iree/compiler/Pipelines/Options.h
+++ b/compiler/src/iree/compiler/Pipelines/Options.h
@@ -92,6 +92,9 @@
// linked into the host program.
// (Currently) only supports the `vmvx-inline` HAL target backend.
InlineStatic = 3,
+ // Inline host-local in-process execution using dynamic executables.
+ // Only supports CPU HAL target backends that produce executable libraries.
+ InlineDynamic = 4,
};
// Program execution model specifying scheduling behavior.
ExecutionModel executionModel = ExecutionModel::AsyncInternal;
diff --git a/compiler/src/iree/compiler/Pipelines/Pipelines.cpp b/compiler/src/iree/compiler/Pipelines/Pipelines.cpp
index 11e3c0e..671d046 100644
--- a/compiler/src/iree/compiler/Pipelines/Pipelines.cpp
+++ b/compiler/src/iree/compiler/Pipelines/Pipelines.cpp
@@ -11,6 +11,7 @@
#include "iree/compiler/Dialect/Flow/Transforms/Passes.h"
#include "iree/compiler/Dialect/HAL/Transforms/Passes.h"
#include "iree/compiler/Dialect/Modules/HAL/Inline/Transforms/Passes.h"
+#include "iree/compiler/Dialect/Modules/HAL/Loader/Transforms/Passes.h"
#include "iree/compiler/Dialect/Stream/Transforms/Passes.h"
#include "iree/compiler/Dialect/Util/Transforms/Passes.h"
#include "iree/compiler/Dialect/VM/Transforms/Passes.h"
@@ -132,6 +133,10 @@
IREE::HAL::Inline::buildHALInlineStaticTransformPassPipeline(
passManager, executableOptions);
break;
+ case SchedulingOptions::ExecutionModel::InlineDynamic:
+ IREE::HAL::Loader::buildHALInlineDynamicTransformPassPipeline(
+ passManager, executableOptions);
+ break;
}
IREE::VM::buildVMTransformPassPipeline(passManager, targetOptions);
diff --git a/compiler/src/iree/compiler/Tools/BUILD b/compiler/src/iree/compiler/Tools/BUILD
index 2bff338..bec1e38 100644
--- a/compiler/src/iree/compiler/Tools/BUILD
+++ b/compiler/src/iree/compiler/Tools/BUILD
@@ -52,6 +52,8 @@
"//compiler/src/iree/compiler/Dialect/HAL/Transforms",
"//compiler/src/iree/compiler/Dialect/Modules/HAL/Inline/IR:HALInlineDialect",
"//compiler/src/iree/compiler/Dialect/Modules/HAL/Inline/Transforms",
+ "//compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/IR:HALLoaderDialect",
+ "//compiler/src/iree/compiler/Dialect/Modules/HAL/Loader/Transforms",
"//compiler/src/iree/compiler/Dialect/Stream/IR",
"//compiler/src/iree/compiler/Dialect/Stream/Transforms",
"//compiler/src/iree/compiler/Dialect/Util/IR",
diff --git a/compiler/src/iree/compiler/Tools/CMakeLists.txt b/compiler/src/iree/compiler/Tools/CMakeLists.txt
index 43da3e5..27b37de 100644
--- a/compiler/src/iree/compiler/Tools/CMakeLists.txt
+++ b/compiler/src/iree/compiler/Tools/CMakeLists.txt
@@ -95,6 +95,8 @@
iree::compiler::Dialect::HAL::Transforms
iree::compiler::Dialect::Modules::HAL::Inline::IR::HALInlineDialect
iree::compiler::Dialect::Modules::HAL::Inline::Transforms
+ iree::compiler::Dialect::Modules::HAL::Loader::IR::HALLoaderDialect
+ iree::compiler::Dialect::Modules::HAL::Loader::Transforms
iree::compiler::Dialect::Stream::IR
iree::compiler::Dialect::Stream::Transforms
iree::compiler::Dialect::Util::IR
diff --git a/compiler/src/iree/compiler/Tools/init_iree_dialects.h b/compiler/src/iree/compiler/Tools/init_iree_dialects.h
index 5cfa88e..9af5377 100644
--- a/compiler/src/iree/compiler/Tools/init_iree_dialects.h
+++ b/compiler/src/iree/compiler/Tools/init_iree_dialects.h
@@ -22,6 +22,7 @@
#include "iree/compiler/Dialect/Flow/IR/FlowDialect.h"
#include "iree/compiler/Dialect/HAL/IR/HALDialect.h"
#include "iree/compiler/Dialect/Modules/HAL/Inline/IR/HALInlineDialect.h"
+#include "iree/compiler/Dialect/Modules/HAL/Loader/IR/HALLoaderDialect.h"
#include "iree/compiler/Dialect/Stream/IR/StreamDialect.h"
#include "iree/compiler/Dialect/Util/IR/UtilDialect.h"
#include "iree/compiler/Dialect/Util/IR/UtilExternalModels.h"
@@ -40,6 +41,7 @@
IREE::Flow::FlowDialect,
IREE::HAL::HALDialect,
IREE::HAL::Inline::HALInlineDialect,
+ IREE::HAL::Loader::HALLoaderDialect,
IREE::LinalgExt::IREELinalgExtDialect,
mlir::linalg::transform::LinalgTransformDialect,
IREE::Stream::StreamDialect,
diff --git a/compiler/src/iree/compiler/Tools/init_iree_passes.h b/compiler/src/iree/compiler/Tools/init_iree_passes.h
index a01be67..443523c 100644
--- a/compiler/src/iree/compiler/Tools/init_iree_passes.h
+++ b/compiler/src/iree/compiler/Tools/init_iree_passes.h
@@ -21,6 +21,7 @@
#include "iree/compiler/Dialect/Flow/Transforms/Passes.h"
#include "iree/compiler/Dialect/HAL/Transforms/Passes.h"
#include "iree/compiler/Dialect/Modules/HAL/Inline/Transforms/Passes.h"
+#include "iree/compiler/Dialect/Modules/HAL/Loader/Transforms/Passes.h"
#include "iree/compiler/Dialect/Stream/Transforms/Passes.h"
#include "iree/compiler/Dialect/Util/Transforms/Passes.h"
#include "iree/compiler/Dialect/VM/Analysis/TestPasses.h"
@@ -48,6 +49,7 @@
IREE::Flow::registerFlowPasses();
IREE::HAL::registerHALPasses();
IREE::HAL::Inline::registerHALInlinePasses();
+ IREE::HAL::Loader::registerHALLoaderPasses();
IREE::LinalgExt::registerPasses();
IREE::Stream::registerStreamPasses();
IREE::Util::registerTransformPasses();
diff --git a/runtime/src/iree/modules/hal/loader/BUILD b/runtime/src/iree/modules/hal/loader/BUILD
new file mode 100644
index 0000000..516935e
--- /dev/null
+++ b/runtime/src/iree/modules/hal/loader/BUILD
@@ -0,0 +1,35 @@
+# Copyright 2022 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
+
+load("//build_tools/bazel:build_defs.oss.bzl", "iree_runtime_cc_library")
+
+package(
+ default_visibility = ["//visibility:public"],
+ features = ["layering_check"],
+ licenses = ["notice"], # Apache 2.0
+)
+
+iree_runtime_cc_library(
+ name = "loader",
+ srcs = [
+ "module.c",
+ ],
+ hdrs = [
+ "module.h",
+ ],
+ textual_hdrs = [
+ "exports.inl",
+ ],
+ deps = [
+ "//runtime/src/iree/base",
+ "//runtime/src/iree/base:tracing",
+ "//runtime/src/iree/hal",
+ "//runtime/src/iree/hal/local:executable_environment",
+ "//runtime/src/iree/hal/local:executable_loader",
+ "//runtime/src/iree/modules/hal:types",
+ "//runtime/src/iree/vm",
+ ],
+)
diff --git a/runtime/src/iree/modules/hal/loader/CMakeLists.txt b/runtime/src/iree/modules/hal/loader/CMakeLists.txt
new file mode 100644
index 0000000..fffb441
--- /dev/null
+++ b/runtime/src/iree/modules/hal/loader/CMakeLists.txt
@@ -0,0 +1,33 @@
+################################################################################
+# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from #
+# runtime/src/iree/modules/hal/loader/BUILD #
+# #
+# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary #
+# CMake-only content. #
+# #
+# To disable autogeneration for this file entirely, delete this header. #
+################################################################################
+
+iree_add_all_subdirs()
+
+iree_cc_library(
+ NAME
+ loader
+ HDRS
+ "module.h"
+ TEXTUAL_HDRS
+ "exports.inl"
+ SRCS
+ "module.c"
+ DEPS
+ iree::base
+ iree::base::tracing
+ iree::hal
+ iree::hal::local::executable_environment
+ iree::hal::local::executable_loader
+ iree::modules::hal::types
+ iree::vm
+ PUBLIC
+)
+
+### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ###
diff --git a/runtime/src/iree/modules/hal/loader/exports.inl b/runtime/src/iree/modules/hal/loader/exports.inl
new file mode 100644
index 0000000..2771c47
--- /dev/null
+++ b/runtime/src/iree/modules/hal/loader/exports.inl
@@ -0,0 +1,31 @@
+// Copyright 2022 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
+
+//===----------------------------------------------------------------------===//
+//
+// ██ ██ █████ ██████ ███ ██ ██ ███ ██ ██████
+// ██ ██ ██ ██ ██ ██ ████ ██ ██ ████ ██ ██
+// ██ █ ██ ███████ ██████ ██ ██ ██ ██ ██ ██ ██ ██ ███
+// ██ ███ ██ ██ ██ ██ ██ ██ ██ ██ ██ ██ ██ ██ ██ ██
+// ███ ███ ██ ██ ██ ██ ██ ████ ██ ██ ████ ██████
+//
+//===----------------------------------------------------------------------===//
+//
+// This file will be auto generated from hal_loader.imports.mlir in the future;
+// for now it's modified by hand but with strict alphabetical sorting required.
+// The order of these functions must be sorted ascending by name in a way
+// compatible with iree_string_view_compare.
+//
+// Users are meant to `#define EXPORT_FN` to be able to access the information.
+// #define EXPORT_FN(name, target_fn, shim_arg_type, arg_type, ret_type)
+
+// clang-format off
+
+EXPORT_FN("executable.dispatch", iree_hal_loader_module_executable_dispatch, dispatch, riiiiCiDCrIID, v)
+EXPORT_FN("executable.load", iree_hal_loader_module_executable_load, rrr, rrr, r)
+EXPORT_FN("executable.query_support", iree_hal_loader_module_executable_query_support, r, r, i)
+
+// clang-format on
diff --git a/runtime/src/iree/modules/hal/loader/module.c b/runtime/src/iree/modules/hal/loader/module.c
new file mode 100644
index 0000000..811fa86
--- /dev/null
+++ b/runtime/src/iree/modules/hal/loader/module.c
@@ -0,0 +1,436 @@
+// Copyright 2022 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 "iree/modules/hal/loader/module.h"
+
+#include "iree/base/api.h"
+#include "iree/base/tracing.h"
+#include "iree/hal/api.h"
+#include "iree/hal/local/local_executable.h"
+#include "iree/vm/api.h"
+
+#define IREE_HAL_LOADER_MODULE_VERSION_0_0 0x00000000u
+#define IREE_HAL_LOADER_MODULE_VERSION_LATEST IREE_HAL_LOADER_MODULE_VERSION_0_0
+
+//===----------------------------------------------------------------------===//
+// Module type definitions
+//===----------------------------------------------------------------------===//
+
+typedef struct iree_hal_loader_module_t {
+ iree_allocator_t host_allocator;
+ iree_hal_loader_module_flags_t flags;
+ // TODO(benvanik): types.
+ iree_host_size_t loader_count;
+ iree_hal_executable_loader_t* loaders[];
+} iree_hal_loader_module_t;
+
+#define IREE_HAL_LOADER_MODULE_CAST(module) \
+ (iree_hal_loader_module_t*)((uint8_t*)(module) + \
+ iree_vm_native_module_size());
+
+typedef struct iree_hal_loader_module_state_t {
+ iree_allocator_t host_allocator;
+ iree_hal_loader_module_flags_t flags;
+} iree_hal_loader_module_state_t;
+
+static void IREE_API_PTR iree_hal_loader_module_destroy(void* base_module) {
+ iree_hal_loader_module_t* module = IREE_HAL_LOADER_MODULE_CAST(base_module);
+ for (iree_host_size_t i = 0; i < module->loader_count; ++i) {
+ iree_hal_executable_loader_release(module->loaders[i]);
+ }
+}
+
+static iree_status_t IREE_API_PTR
+iree_hal_loader_module_alloc_state(void* self, iree_allocator_t host_allocator,
+ iree_vm_module_state_t** out_module_state) {
+ IREE_TRACE_ZONE_BEGIN(z0);
+
+ iree_hal_loader_module_t* module = IREE_HAL_LOADER_MODULE_CAST(self);
+ iree_hal_loader_module_state_t* state = NULL;
+ IREE_RETURN_AND_END_ZONE_IF_ERROR(
+ z0,
+ iree_allocator_malloc(host_allocator, sizeof(*state), (void**)&state));
+ memset(state, 0, sizeof(*state));
+ state->host_allocator = host_allocator;
+ state->flags = module->flags;
+
+ *out_module_state = (iree_vm_module_state_t*)state;
+ IREE_TRACE_ZONE_END(z0);
+ return iree_ok_status();
+}
+
+static void IREE_API_PTR iree_hal_loader_module_free_state(
+ void* self, iree_vm_module_state_t* module_state) {
+ IREE_TRACE_ZONE_BEGIN(z0);
+
+ iree_hal_loader_module_state_t* state =
+ (iree_hal_loader_module_state_t*)module_state;
+ iree_allocator_free(state->host_allocator, state);
+
+ IREE_TRACE_ZONE_END(z0);
+}
+
+static iree_status_t IREE_API_PTR iree_hal_loader_module_notify(
+ void* self, iree_vm_module_state_t* module_state, iree_vm_signal_t signal) {
+ switch (signal) {
+ case IREE_VM_SIGNAL_SUSPEND:
+ case IREE_VM_SIGNAL_LOW_MEMORY:
+ default:
+ return iree_ok_status();
+ }
+}
+
+//===----------------------------------------------------------------------===//
+// Utilities
+//===----------------------------------------------------------------------===//
+
+// Casts a VM value to a C host size.
+static iree_host_size_t iree_hal_cast_host_size(int64_t value) {
+ // TODO(benvanik): make this return status and check for overflow if host
+ // size is 32-bits.
+ return (iree_host_size_t)value;
+}
+
+// Casts a VM value to a HAL device size.
+static iree_device_size_t iree_hal_cast_device_size(int64_t value) {
+ // TODO(benvanik): make this return status and check for overflow if device
+ // size is 32-bits.
+ return (iree_device_size_t)value;
+}
+
+//===----------------------------------------------------------------------===//
+// Shared argument shims
+//===----------------------------------------------------------------------===//
+
+#define IREE_HAL_ABI_EXPORT(function_name, arg_types, ret_types) \
+ IREE_VM_ABI_EXPORT(function_name, iree_hal_loader_module_state_t, arg_types, \
+ ret_types)
+#define IREE_HAL_ABI_FIXED_STRUCT(name, types, body) \
+ IREE_VM_ABI_FIXED_STRUCT(name, body)
+#define IREE_HAL_ABI_DEFINE_SHIM(arg_types, ret_types) \
+ static IREE_VM_ABI_DEFINE_SHIM(arg_types, ret_types)
+
+//===----------------------------------------------------------------------===//
+// iree_hal_executable_t
+//===----------------------------------------------------------------------===//
+
+IREE_HAL_ABI_EXPORT(iree_hal_loader_module_executable_query_support, //
+ r, i) {
+ iree_vm_buffer_t* executable_format = NULL;
+ IREE_RETURN_IF_ERROR(
+ iree_vm_buffer_check_deref(args->r0, &executable_format));
+ iree_string_view_t executable_format_str =
+ iree_vm_buffer_as_string(executable_format);
+
+ bool has_support = false;
+ iree_hal_loader_module_t* loader_module = IREE_HAL_LOADER_MODULE_CAST(module);
+ for (iree_host_size_t i = 0; i < loader_module->loader_count; ++i) {
+ iree_hal_executable_loader_t* loader = loader_module->loaders[i];
+ if (iree_hal_executable_loader_query_support(loader, 0,
+ executable_format_str)) {
+ has_support = true;
+ break;
+ }
+ }
+
+ rets->i0 = has_support ? 1 : 0;
+ return iree_ok_status();
+}
+
+static iree_status_t iree_hal_loader_module_try_load(
+ iree_hal_loader_module_t* loader_module,
+ const iree_hal_executable_params_t* executable_params,
+ iree_hal_executable_t** out_executable) {
+ for (iree_host_size_t i = 0; i < loader_module->loader_count; ++i) {
+ iree_hal_executable_loader_t* loader = loader_module->loaders[i];
+ if (!iree_hal_executable_loader_query_support(
+ loader, executable_params->caching_mode,
+ executable_params->executable_format)) {
+ // Loader definitely can't handle the executable; no use trying so skip.
+ continue;
+ }
+ // The loader _may_ handle the executable; if the specific executable is not
+ // supported then the try will fail with IREE_STATUS_CANCELLED and we should
+ // continue trying other loaders.
+ iree_status_t status = iree_hal_executable_loader_try_load(
+ loader, executable_params, out_executable);
+ if (iree_status_is_ok(status)) {
+ // Executable was successfully loaded.
+ return status;
+ } else if (!iree_status_is_cancelled(status)) {
+ // Error beyond just the try failing due to unsupported formats.
+ return status;
+ }
+ iree_status_ignore(status);
+ }
+ return iree_make_status(
+ IREE_STATUS_NOT_FOUND,
+ "no executable loader registered for the given executable format '%.*s'",
+ (int)executable_params->executable_format.size,
+ executable_params->executable_format.data);
+}
+
+IREE_HAL_ABI_EXPORT(iree_hal_loader_module_executable_load, //
+ rrr, r) {
+ iree_vm_buffer_t* executable_format = NULL;
+ IREE_RETURN_IF_ERROR(
+ iree_vm_buffer_check_deref(args->r0, &executable_format));
+ iree_string_view_t executable_format_str =
+ iree_vm_buffer_as_string(executable_format);
+ iree_vm_buffer_t* executable_data = NULL;
+ IREE_RETURN_IF_ERROR(iree_vm_buffer_check_deref(args->r1, &executable_data));
+ iree_host_size_t constant_count = 0;
+ const uint32_t* constants = NULL;
+ if (iree_vm_buffer_isa(args->r2)) {
+ iree_vm_buffer_t* constant_buffer = NULL;
+ IREE_RETURN_IF_ERROR(
+ iree_vm_buffer_check_deref(args->r2, &constant_buffer));
+ if (constant_buffer->data.data_length % 4 != 0) {
+ return iree_make_status(IREE_STATUS_INVALID_ARGUMENT,
+ "constant buffer data must contain 4-byte "
+ "elements but data length is %" PRIhsz,
+ constant_buffer->data.data_length);
+ }
+ constant_count = constant_buffer->data.data_length / sizeof(uint32_t);
+ constants = (const uint32_t*)constant_buffer->data.data;
+ }
+
+ iree_hal_executable_params_t executable_params;
+ iree_hal_executable_params_initialize(&executable_params);
+ executable_params.caching_mode |=
+ executable_data->access == IREE_VM_BUFFER_ACCESS_ORIGIN_MODULE
+ ? IREE_HAL_EXECUTABLE_CACHING_MODE_ALIAS_PROVIDED_DATA
+ : 0;
+ executable_params.executable_format = executable_format_str;
+ executable_params.executable_data = iree_make_const_byte_span(
+ executable_data->data.data, executable_data->data.data_length);
+ executable_params.executable_layout_count = 0;
+ executable_params.executable_layouts = NULL;
+ executable_params.constant_count = constant_count;
+ executable_params.constants = constants;
+
+ iree_hal_executable_t* executable = NULL;
+ iree_hal_loader_module_t* loader_module = IREE_HAL_LOADER_MODULE_CAST(module);
+ iree_status_t status = iree_hal_loader_module_try_load(
+ loader_module, &executable_params, &executable);
+
+ rets->r0 = iree_hal_executable_move_ref(executable);
+ return status;
+}
+
+typedef struct {
+ union {
+ struct {
+ iree_vm_ref_t executable;
+ int32_t entry_point;
+ int32_t workgroup_x;
+ int32_t workgroup_y;
+ int32_t workgroup_z;
+ };
+ iree_vm_abi_riiii_t params;
+ };
+ iree_vm_size_t push_constant_count;
+ const uint32_t* push_constants;
+ iree_vm_size_t binding_count;
+ const iree_vm_abi_rII_t* bindings;
+} iree_hal_loader_dispatch_args_t;
+
+static iree_status_t iree_hal_loader_module_executable_dispatch(
+ iree_vm_stack_t* IREE_RESTRICT stack, void* IREE_RESTRICT module,
+ iree_hal_loader_module_state_t* IREE_RESTRICT state,
+ const iree_hal_loader_dispatch_args_t* IREE_RESTRICT args) {
+ iree_hal_executable_t* executable = NULL;
+ IREE_RETURN_IF_ERROR(
+ iree_hal_executable_check_deref(args->executable, &executable));
+
+ if (args->binding_count > 32) {
+ return iree_make_status(IREE_STATUS_RESOURCE_EXHAUSTED,
+ "too many bindings");
+ }
+ void** binding_ptrs =
+ (void**)iree_alloca(args->binding_count * sizeof(void*));
+ size_t* binding_lengths =
+ (size_t*)iree_alloca(args->binding_count * sizeof(size_t));
+ for (iree_host_size_t i = 0; i < args->binding_count; ++i) {
+ iree_vm_buffer_t* buffer = NULL;
+ IREE_RETURN_IF_ERROR(
+ iree_vm_buffer_check_deref(args->bindings[i].r0, &buffer));
+ // TODO(benvanik): this is a hack around not having the access permissions
+ // currently modeled. This is only used for verification and early errors
+ // and not intended to be a last-line defense against writes (you need an
+ // MMU for that) so it's just subpar reporting.
+ iree_const_byte_span_t span;
+ IREE_RETURN_IF_ERROR(iree_vm_buffer_map_ro(
+ buffer, iree_hal_cast_host_size(args->bindings[i].i1),
+ iree_hal_cast_host_size(args->bindings[i].i2), /*alignment=*/1, &span));
+ binding_ptrs[i] = (void*)span.data;
+ binding_lengths[i] = span.data_length;
+ }
+
+ const iree_hal_executable_dispatch_state_v0_t dispatch_state = {
+ .workgroup_size_x = 1,
+ .workgroup_size_y = 1,
+ .workgroup_size_z = 1,
+ .push_constant_count = args->push_constant_count,
+ .workgroup_count_x = args->workgroup_x,
+ .workgroup_count_y = args->workgroup_y,
+ .workgroup_count_z = args->workgroup_z,
+ .max_concurrency = 1,
+ .binding_count = args->binding_count,
+ .push_constants = args->push_constants,
+ .binding_ptrs = binding_ptrs,
+ .binding_lengths = binding_lengths,
+ };
+
+ // TODO(benvanik): environmental information.
+ uint32_t processor_id = 0;
+ iree_byte_span_t local_memory = iree_byte_span_empty();
+
+ return iree_hal_local_executable_issue_dispatch_inline(
+ (iree_hal_local_executable_t*)executable, args->entry_point,
+ &dispatch_state, processor_id, local_memory);
+}
+
+static iree_status_t iree_vm_shim_dispatch_v(
+ iree_vm_stack_t* IREE_RESTRICT stack, iree_vm_native_function_flags_t flags,
+ iree_byte_span_t args_storage, iree_byte_span_t rets_storage,
+ iree_vm_native_function_target2_t target_fn, void* IREE_RESTRICT module,
+ void* IREE_RESTRICT module_state) {
+ // TODO(benvanik): support multiple variadic segments in one call.
+ // For now we inline what it would do in a very painful way.
+ bool args_ok = true;
+ if (args_storage.data_length <
+ (sizeof(iree_vm_abi_riiii_t) + sizeof(iree_vm_size_t) +
+ sizeof(iree_vm_size_t))) {
+ // Can't fit even with zero lengths.
+ args_ok = false;
+ }
+ iree_hal_loader_dispatch_args_t args = {
+ .params = *(const iree_vm_abi_riiii_t*)args_storage.data,
+ };
+ if (args_ok) {
+ const uint8_t* push_constants_ptr = args_storage.data + sizeof(args.params);
+ args.push_constant_count = *(const iree_vm_size_t*)push_constants_ptr;
+ args.push_constants =
+ (const uint32_t*)(push_constants_ptr + sizeof(iree_vm_size_t));
+ const uint8_t* bindings_ptr =
+ push_constants_ptr + sizeof(iree_vm_size_t) +
+ args.push_constant_count * sizeof(args.push_constants[0]);
+ args.binding_count = *(const iree_vm_size_t*)bindings_ptr;
+ args.bindings =
+ (const iree_vm_abi_rII_t*)(bindings_ptr + sizeof(iree_vm_size_t));
+ const uint8_t* max_ptr = (const uint8_t*)args.bindings +
+ args.binding_count * sizeof(args.bindings[0]);
+ const uint8_t* end_ptr = args_storage.data + args_storage.data_length;
+ if (max_ptr > end_ptr) args_ok = false;
+ }
+ if (IREE_UNLIKELY(!args_ok || rets_storage.data_length > 0)) {
+ return iree_make_status(IREE_STATUS_INVALID_ARGUMENT,
+ "argument/result signature mismatch");
+ }
+ return iree_hal_loader_module_executable_dispatch(stack, module, module_state,
+ &args);
+}
+
+//===----------------------------------------------------------------------===//
+// VM module interface implementation
+//===----------------------------------------------------------------------===//
+
+// NOTE: this must match the ordering of the iree_hal_loader_module_exports_
+// table.
+static const iree_vm_native_function_ptr_t iree_hal_loader_module_funcs_[] = {
+#define EXPORT_FN(name, target_fn, shim_arg_type, arg_types, ret_types) \
+ { \
+ .shim = (iree_vm_native_function_shim_t) \
+ iree_vm_shim_##shim_arg_type##_##ret_types, \
+ .target = (iree_vm_native_function_target_t)(target_fn), \
+ },
+#include "iree/modules/hal/loader/exports.inl" // IWYU pragma: keep
+#undef EXPORT_FN
+};
+
+// NOTE: 0 length, but can't express that in C.
+static const iree_vm_native_import_descriptor_t
+ iree_hal_loader_module_imports_[1];
+
+static const iree_vm_native_export_descriptor_t
+ iree_hal_loader_module_exports_[] = {
+#define EXPORT_FN(name, target_fn, shim_arg_type, arg_types, ret_types) \
+ { \
+ .local_name = iree_string_view_literal(name), \
+ .calling_convention = \
+ iree_string_view_literal("0" #arg_types "_" #ret_types), \
+ .attr_count = 0, \
+ .attrs = NULL, \
+ },
+#include "iree/modules/hal/loader/exports.inl" // IWYU pragma: keep
+#undef EXPORT_FN
+};
+static_assert(IREE_ARRAYSIZE(iree_hal_loader_module_funcs_) ==
+ IREE_ARRAYSIZE(iree_hal_loader_module_exports_),
+ "function pointer table must be 1:1 with exports");
+
+static const iree_vm_native_module_descriptor_t
+ iree_hal_loader_module_descriptor_ = {
+ .name = iree_string_view_literal("hal_loader"),
+ .version = IREE_HAL_LOADER_MODULE_VERSION_LATEST,
+ .attr_count = 0,
+ .attrs = NULL,
+ .dependency_count = 0,
+ .dependencies = NULL,
+ .import_count = 0, // workaround for 0-length C struct
+ .imports = iree_hal_loader_module_imports_,
+ .export_count = IREE_ARRAYSIZE(iree_hal_loader_module_exports_),
+ .exports = iree_hal_loader_module_exports_,
+ .function_count = IREE_ARRAYSIZE(iree_hal_loader_module_funcs_),
+ .functions = iree_hal_loader_module_funcs_,
+};
+
+IREE_API_EXPORT iree_status_t iree_hal_loader_module_create(
+ iree_vm_instance_t* instance, iree_hal_loader_module_flags_t flags,
+ iree_host_size_t loader_count, iree_hal_executable_loader_t** loaders,
+ iree_allocator_t host_allocator, iree_vm_module_t** out_module) {
+ IREE_ASSERT_ARGUMENT(instance);
+ IREE_ASSERT_ARGUMENT(out_module);
+ *out_module = NULL;
+
+ // Setup the interface with the functions we implement ourselves. Any function
+ // we omit will be handled by the base native module.
+ static const iree_vm_module_t interface = {
+ .destroy = iree_hal_loader_module_destroy,
+ .alloc_state = iree_hal_loader_module_alloc_state,
+ .free_state = iree_hal_loader_module_free_state,
+ .notify = iree_hal_loader_module_notify,
+ };
+
+ // Allocate shared module state.
+ iree_host_size_t total_size =
+ iree_vm_native_module_size() + sizeof(iree_hal_loader_module_t) +
+ loader_count * sizeof(iree_hal_executable_loader_t*);
+ iree_vm_module_t* base_module = NULL;
+ IREE_RETURN_IF_ERROR(
+ iree_allocator_malloc(host_allocator, total_size, (void**)&base_module));
+ memset(base_module, 0, total_size);
+ iree_status_t status = iree_vm_native_module_initialize(
+ &interface, &iree_hal_loader_module_descriptor_, instance, host_allocator,
+ base_module);
+ if (!iree_status_is_ok(status)) {
+ iree_allocator_free(host_allocator, base_module);
+ return status;
+ }
+
+ iree_hal_loader_module_t* module = IREE_HAL_LOADER_MODULE_CAST(base_module);
+ module->host_allocator = host_allocator;
+ module->flags = flags;
+ module->loader_count = loader_count;
+ for (iree_host_size_t i = 0; i < loader_count; ++i) {
+ module->loaders[i] = loaders[i];
+ iree_hal_executable_loader_retain(loaders[i]);
+ }
+
+ *out_module = base_module;
+ return iree_ok_status();
+}
diff --git a/runtime/src/iree/modules/hal/loader/module.h b/runtime/src/iree/modules/hal/loader/module.h
new file mode 100644
index 0000000..1078acd
--- /dev/null
+++ b/runtime/src/iree/modules/hal/loader/module.h
@@ -0,0 +1,37 @@
+// Copyright 2022 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
+
+#ifndef IREE_MODULES_HAL_LOADER_MODULE_H_
+#define IREE_MODULES_HAL_LOADER_MODULE_H_
+
+#include <stdint.h>
+
+#include "iree/base/api.h"
+#include "iree/hal/api.h"
+#include "iree/hal/local/executable_loader.h"
+#include "iree/modules/hal/types.h"
+#include "iree/vm/api.h"
+
+#ifdef __cplusplus
+extern "C" {
+#endif // __cplusplus
+
+enum iree_hal_loader_module_flag_bits_t {
+ IREE_HAL_LOADER_MODULE_FLAG_NONE = 0u,
+};
+typedef uint32_t iree_hal_loader_module_flags_t;
+
+// Creates the dynamic HAL executable loader module for local execution.
+IREE_API_EXPORT iree_status_t iree_hal_loader_module_create(
+ iree_vm_instance_t* instance, iree_hal_loader_module_flags_t flags,
+ iree_host_size_t loader_count, iree_hal_executable_loader_t** loaders,
+ iree_allocator_t host_allocator, iree_vm_module_t** out_module);
+
+#ifdef __cplusplus
+} // extern "C"
+#endif // __cplusplus
+
+#endif // IREE_MODULES_HAL_LOADER_MODULE_H_
diff --git a/runtime/src/iree/tooling/BUILD b/runtime/src/iree/tooling/BUILD
index 605c1fb..1992b2f 100644
--- a/runtime/src/iree/tooling/BUILD
+++ b/runtime/src/iree/tooling/BUILD
@@ -23,8 +23,11 @@
"//runtime/src/iree/base/internal:file_io",
"//runtime/src/iree/base/internal:flags",
"//runtime/src/iree/hal",
+ "//runtime/src/iree/hal/local/loaders/registration",
"//runtime/src/iree/modules/hal",
"//runtime/src/iree/modules/hal/inline",
+ "//runtime/src/iree/modules/hal/loader",
+ "//runtime/src/iree/modules/vmvx",
"//runtime/src/iree/vm",
"//runtime/src/iree/vm:bytecode_module",
],
diff --git a/runtime/src/iree/tooling/CMakeLists.txt b/runtime/src/iree/tooling/CMakeLists.txt
index f36aee9..552d6ab 100644
--- a/runtime/src/iree/tooling/CMakeLists.txt
+++ b/runtime/src/iree/tooling/CMakeLists.txt
@@ -24,8 +24,11 @@
iree::base::internal::flags
iree::base::tracing
iree::hal
+ iree::hal::local::loaders::registration
iree::modules::hal
iree::modules::hal::inline
+ iree::modules::hal::loader
+ iree::modules::vmvx
iree::vm
iree::vm::bytecode_module
PUBLIC
diff --git a/runtime/src/iree/tooling/context_util.c b/runtime/src/iree/tooling/context_util.c
index 85e4d7d..988bf3b 100644
--- a/runtime/src/iree/tooling/context_util.c
+++ b/runtime/src/iree/tooling/context_util.c
@@ -13,7 +13,9 @@
#include "iree/base/internal/file_io.h"
#include "iree/base/internal/flags.h"
#include "iree/base/tracing.h"
+#include "iree/hal/local/loaders/registration/init.h"
#include "iree/modules/hal/inline/module.h"
+#include "iree/modules/hal/loader/module.h"
#include "iree/modules/hal/module.h"
#include "iree/tooling/device_util.h"
#include "iree/vm/bytecode_module.h"
@@ -191,6 +193,47 @@
return status;
}
+static iree_status_t iree_tooling_load_hal_loader_module(
+ iree_vm_instance_t* instance, iree_allocator_t host_allocator,
+ iree_vm_module_t** out_module) {
+ IREE_ASSERT_ARGUMENT(instance);
+ IREE_ASSERT_ARGUMENT(out_module);
+ *out_module = NULL;
+ IREE_TRACE_ZONE_BEGIN(z0);
+
+ // Register required types before creating the module.
+ IREE_RETURN_AND_END_ZONE_IF_ERROR(
+ z0, iree_hal_module_register_loader_types(instance));
+
+ // Create all executable loaders built into the binary.
+ // We could allow users to choose the set with a flag.
+ iree_host_size_t loader_count = 0;
+ iree_hal_executable_loader_t* loaders[16];
+ iree_status_t status = iree_hal_create_all_available_executable_loaders(
+ IREE_ARRAYSIZE(loaders), &loader_count, loaders, host_allocator);
+
+ // Create the module; it retains the loaders for its lifetime.
+ iree_vm_module_t* module = NULL;
+ if (iree_status_is_ok(status)) {
+ iree_hal_loader_module_flags_t flags = IREE_HAL_LOADER_MODULE_FLAG_NONE;
+ status = iree_hal_loader_module_create(instance, flags, loader_count,
+ loaders, host_allocator, &module);
+ }
+
+ // Always release loaders; loader module has retained them.
+ for (iree_host_size_t i = 0; i < loader_count; ++i) {
+ iree_hal_executable_loader_release(loaders[i]);
+ }
+
+ if (iree_status_is_ok(status)) {
+ *out_module = module;
+ } else {
+ iree_vm_module_release(module);
+ }
+ IREE_TRACE_ZONE_END(z0);
+ return status;
+}
+
//===----------------------------------------------------------------------===//
// Module management
//===----------------------------------------------------------------------===//
@@ -275,6 +318,9 @@
IREE_RETURN_IF_ERROR(iree_tooling_load_hal_inline_module(
state->instance, state->host_allocator, &module,
&state->device_allocator));
+ } else if (iree_string_view_equal(dependency->name, IREE_SV("hal_loader"))) {
+ IREE_RETURN_IF_ERROR(iree_tooling_load_hal_loader_module(
+ state->instance, state->host_allocator, &module));
} else if (iree_string_view_equal(dependency->name, IREE_SV("vmvx"))) {
IREE_RETURN_IF_ERROR(iree_vmvx_module_create(
state->instance, state->host_allocator, &module));
diff --git a/runtime/src/iree/vm/shims.c b/runtime/src/iree/vm/shims.c
index 57707d7..e0890c0 100644
--- a/runtime/src/iree/vm/shims.c
+++ b/runtime/src/iree/vm/shims.c
@@ -48,6 +48,7 @@
IREE_VM_ABI_DEFINE_SHIM(rr, ii);
IREE_VM_ABI_DEFINE_SHIM(rr, iI);
IREE_VM_ABI_DEFINE_SHIM(rrr, iI);
+IREE_VM_ABI_DEFINE_SHIM(rrr, r);
IREE_VM_ABI_DEFINE_SHIM(rrCirIID, r);
IREE_VM_ABI_DEFINE_SHIM(rriCiD, v);
IREE_VM_ABI_DEFINE_SHIM(rriiCID, v);
diff --git a/runtime/src/iree/vm/shims.h b/runtime/src/iree/vm/shims.h
index f7f3f38..0682329 100644
--- a/runtime/src/iree/vm/shims.h
+++ b/runtime/src/iree/vm/shims.h
@@ -285,6 +285,14 @@
int32_t i3;
});
+IREE_VM_ABI_FIXED_STRUCT(riiii, {
+ iree_vm_ref_t r0;
+ int32_t i1;
+ int32_t i2;
+ int32_t i3;
+ int32_t i4;
+});
+
IREE_VM_ABI_FIXED_STRUCT(riiI, {
iree_vm_ref_t r0;
int32_t i1;
@@ -564,6 +572,7 @@
IREE_VM_ABI_DECLARE_SHIM(rr, ii);
IREE_VM_ABI_DECLARE_SHIM(rr, iI);
IREE_VM_ABI_DECLARE_SHIM(rrr, iI);
+IREE_VM_ABI_DECLARE_SHIM(rrr, r);
IREE_VM_ABI_DECLARE_SHIM(rrCirIID, r);
IREE_VM_ABI_DECLARE_SHIM(rriCiD, v);
IREE_VM_ABI_DECLARE_SHIM(rriiCID, v);
