Delete llvmjit compiler target and HAL backend. (#4152)
LLVM AOT is now stable enough to run by default on all build/test configurations (mainly thanks to: https://github.com/google/iree/pull/3757, https://github.com/google/iree/pull/4034, https://github.com/google/iree/pull/3928).
We don't see a path to deployment on our primary platforms using LLVM JIT and the backend adds excess maintenance cost in the main repository (e.g. runtime code that depends on LLVM -> slow builds and can't build on Android), so this removes the `llvm-ir` compiler target and corresponding `llvm` HAL backend in favor of the `dylib-llvm-aot` compiler target and `dylib` HAL backend (we have many names for the same components :P).
diff --git a/CMakeLists.txt b/CMakeLists.txt
index d72050b..cc96638 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -106,7 +106,6 @@
# List of all HAL drivers to be built by default:
set(IREE_ALL_HAL_DRIVERS
DyLib
- LLVM
Metal
VMLA
Vulkan
@@ -114,10 +113,6 @@
if(IREE_HAL_DRIVERS_TO_BUILD STREQUAL "all")
set(IREE_HAL_DRIVERS_TO_BUILD ${IREE_ALL_HAL_DRIVERS})
- # For cross compilation towards Android, we don't want LLVM JIT HAL driver.
- if(ANDROID)
- list(REMOVE_ITEM IREE_HAL_DRIVERS_TO_BUILD LLVM)
- endif()
# For Apple platforms we need to use Metal instead of Vulkan.
if(APPLE)
@@ -144,7 +139,6 @@
# List of all target backends to be built by default:
set(IREE_ALL_TARGET_BACKENDS
DYLIB-LLVM-AOT
- LLVM-IR
Metal-SPIRV
Vulkan-SPIRV
VMLA
diff --git a/docs/get_started/getting_started_linux_cmake.md b/docs/get_started/getting_started_linux_cmake.md
index 83d0fb2..502123b 100644
--- a/docs/get_started/getting_started_linux_cmake.md
+++ b/docs/get_started/getting_started_linux_cmake.md
@@ -109,7 +109,7 @@
### LLVM Ahead-of-Time (AOT) backend
-To compile IREE LLVM AOT (vs JIT) module we need to set the AOT linker path environment variable:
+To compile IREE LLVM AOT module we need to set the AOT linker path environment variable:
```shell
$ export IREE_LLVMAOT_LINKER_PATH=ld.lld-10
diff --git a/docs/tensorflow_coverage/language_and_speech_coverage.md b/docs/tensorflow_coverage/language_and_speech_coverage.md
index f49581e..789a8dc 100644
--- a/docs/tensorflow_coverage/language_and_speech_coverage.md
+++ b/docs/tensorflow_coverage/language_and_speech_coverage.md
@@ -2,9 +2,9 @@
Tests of MobileBert and streamable Keyword Spotting models.
-IREE has three backend
+IREE has three main backend
[targets](https://github.com/google/iree/tree/main/iree/compiler/Dialect/HAL/Target):
-`vmla` , `llvm-ir` and `vulkan-spirv`. We also test TFLite in our
+`vmla` , `llvm` and `vulkan-spirv`. We also test TFLite in our
infrastructure for benchmarking purposes.
*Last Updated: 2020/12/8*
diff --git a/docs/tensorflow_coverage/tf_base_coverage.md b/docs/tensorflow_coverage/tf_base_coverage.md
index f2ccda3..4fb1fce 100644
--- a/docs/tensorflow_coverage/tf_base_coverage.md
+++ b/docs/tensorflow_coverage/tf_base_coverage.md
@@ -2,9 +2,9 @@
Tests of the `tf` , `tf.math` , `tf.nn` , `tf.signal` and `tf.strings` APIs.
-IREE has three backend
+IREE has three main backend
[targets](https://github.com/google/iree/tree/main/iree/compiler/Dialect/HAL/Target):
-`vmla` , `llvm-ir` and `vulkan-spirv`. We also test TFLite in our
+`vmla` , `llvm` and `vulkan-spirv`. We also test TFLite in our
infrastructure for benchmarking purposes.
*Last Updated: 2020/12/8*
diff --git a/docs/tensorflow_coverage/tf_keras_coverage.md b/docs/tensorflow_coverage/tf_keras_coverage.md
index 97a962a..325c300 100644
--- a/docs/tensorflow_coverage/tf_keras_coverage.md
+++ b/docs/tensorflow_coverage/tf_keras_coverage.md
@@ -2,9 +2,9 @@
Tests of `tf.keras.layers` compiled with static shapes, dynamic shapes and training enabled.
-IREE has three backend
+IREE has three main backend
[targets](https://github.com/google/iree/tree/main/iree/compiler/Dialect/HAL/Target):
-`vmla` , `llvm-ir` and `vulkan-spirv`. We also test TFLite in our
+`vmla` , `llvm` and `vulkan-spirv`. We also test TFLite in our
infrastructure for benchmarking purposes.
*Last Updated: 2020/12/8*
diff --git a/docs/tensorflow_coverage/vision_coverage.md b/docs/tensorflow_coverage/vision_coverage.md
index 1105714..e543434 100644
--- a/docs/tensorflow_coverage/vision_coverage.md
+++ b/docs/tensorflow_coverage/vision_coverage.md
@@ -2,9 +2,9 @@
Tests of Keras and Slim vision models.
-IREE has three backend
+IREE has three main backend
[targets](https://github.com/google/iree/tree/main/iree/compiler/Dialect/HAL/Target):
-`vmla` , `llvm-ir` and `vulkan-spirv`. We also test TFLite in our
+`vmla` , `llvm` and `vulkan-spirv`. We also test TFLite in our
infrastructure for benchmarking purposes.
*Last Updated: 2020/12/8*
diff --git a/iree/compiler/Conversion/LinalgToLLVM/ConvertToLLVM.cpp b/iree/compiler/Conversion/LinalgToLLVM/ConvertToLLVM.cpp
index 7a2c0eb..2afcd3b 100644
--- a/iree/compiler/Conversion/LinalgToLLVM/ConvertToLLVM.cpp
+++ b/iree/compiler/Conversion/LinalgToLLVM/ConvertToLLVM.cpp
@@ -438,8 +438,7 @@
static PassRegistration<ConvertToLLVMPass> pass(
"iree-codegen-convert-to-llvm",
"Perform final conversion from Linalg/HAL/Shape/Vector/Standard to "
- "LLVMIR "
- "dialect",
+ "LLVMIR dialect",
[] { return std::make_unique<ConvertToLLVMPass>(); });
} // namespace iree_compiler
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/BUILD b/iree/compiler/Dialect/HAL/Target/LLVM/BUILD
index 4f29ec2..cd077ec 100644
--- a/iree/compiler/Dialect/HAL/Target/LLVM/BUILD
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/BUILD
@@ -22,7 +22,7 @@
iree_cmake_extra_content(
content = """
-if(NOT ${IREE_TARGET_BACKEND_LLVM-IR} AND NOT ${IREE_TARGET_BACKEND_DYLIB-LLVM-AOT})
+if(NOT ${IREE_TARGET_BACKEND_DYLIB-LLVM-AOT})
return()
endif()
""",
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/CMakeLists.txt b/iree/compiler/Dialect/HAL/Target/LLVM/CMakeLists.txt
index a0768ba..b968fe5 100644
--- a/iree/compiler/Dialect/HAL/Target/LLVM/CMakeLists.txt
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/CMakeLists.txt
@@ -12,7 +12,7 @@
# See the License for the specific language governing permissions and
# limitations under the License.
-if(NOT ${IREE_TARGET_BACKEND_LLVM-IR} AND NOT ${IREE_TARGET_BACKEND_DYLIB-LLVM-AOT})
+if(NOT ${IREE_TARGET_BACKEND_DYLIB-LLVM-AOT})
return()
endif()
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/IR/BUILD b/iree/compiler/Dialect/HAL/Target/LLVM/IR/BUILD
deleted file mode 100644
index 10b20b5..0000000
--- a/iree/compiler/Dialect/HAL/Target/LLVM/IR/BUILD
+++ /dev/null
@@ -1,52 +0,0 @@
-# Copyright 2020 Google LLC
-#
-# Licensed under the Apache License, Version 2.0 (the "License");
-# you may not use this file except in compliance with the License.
-# You may obtain a copy of the License at
-#
-# https://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing, software
-# distributed under the License is distributed on an "AS IS" BASIS,
-# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-# See the License for the specific language governing permissions and
-# limitations under the License.
-
-load("//iree:build_defs.oss.bzl", "iree_cmake_extra_content")
-
-package(
- default_visibility = ["//visibility:public"],
- features = ["layering_check"],
- licenses = ["notice"], # Apache 2.0
-)
-
-iree_cmake_extra_content(
- content = """
-if(NOT ${IREE_TARGET_BACKEND_LLVM-IR})
- return()
-endif()
-""",
-)
-
-cc_library(
- name = "LLVMIR",
- srcs = [
- "LLVMIRTarget.cpp",
- ],
- hdrs = [
- "LLVMIRTarget.h",
- ],
- deps = [
- "//iree/base:flatcc",
- "//iree/compiler/Dialect/HAL/Target",
- "//iree/compiler/Dialect/HAL/Target/LLVM:LLVMBaseTarget",
- "//iree/compiler/Dialect/HAL/Target/LLVM:LLVMIRPasses",
- "//iree/compiler/Dialect/HAL/Target/LLVM:LLVMTargetOptions",
- "//iree/compiler/Dialect/Shape/IR",
- "//iree/compiler/Utils",
- "//iree/schemas:llvmir_executable_def_c_fbs",
- "@llvm-project//llvm:Core",
- "@llvm-project//llvm:Support",
- "@llvm-project//mlir:TargetLLVMIR",
- ],
-)
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/IR/CMakeLists.txt b/iree/compiler/Dialect/HAL/Target/LLVM/IR/CMakeLists.txt
deleted file mode 100644
index 650ce95..0000000
--- a/iree/compiler/Dialect/HAL/Target/LLVM/IR/CMakeLists.txt
+++ /dev/null
@@ -1,41 +0,0 @@
-# Copyright 2020 Google LLC
-#
-# Licensed under the Apache License, Version 2.0 (the "License");
-# you may not use this file except in compliance with the License.
-# You may obtain a copy of the License at
-#
-# https://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing, software
-# distributed under the License is distributed on an "AS IS" BASIS,
-# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-# See the License for the specific language governing permissions and
-# limitations under the License.
-
-if(NOT ${IREE_TARGET_BACKEND_LLVM-IR})
- return()
-endif()
-
-iree_add_all_subdirs()
-
-iree_cc_library(
- NAME
- LLVMIR
- HDRS
- "LLVMIRTarget.h"
- SRCS
- "LLVMIRTarget.cpp"
- DEPS
- LLVMCore
- LLVMSupport
- MLIRTargetLLVMIR
- iree::base::flatcc
- iree::compiler::Dialect::HAL::Target
- iree::compiler::Dialect::HAL::Target::LLVM::LLVMBaseTarget
- iree::compiler::Dialect::HAL::Target::LLVM::LLVMIRPasses
- iree::compiler::Dialect::HAL::Target::LLVM::LLVMTargetOptions
- iree::compiler::Dialect::Shape::IR
- iree::compiler::Utils
- iree::schemas::llvmir_executable_def_c_fbs
- PUBLIC
-)
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/IR/LLVMIRTarget.cpp b/iree/compiler/Dialect/HAL/Target/LLVM/IR/LLVMIRTarget.cpp
deleted file mode 100644
index 93f5d20..0000000
--- a/iree/compiler/Dialect/HAL/Target/LLVM/IR/LLVMIRTarget.cpp
+++ /dev/null
@@ -1,119 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#include "iree/compiler/Dialect/HAL/Target/LLVM/IR/LLVMIRTarget.h"
-
-#include "iree/compiler/Dialect/HAL/Target/LLVM/LLVMBaseTarget.h"
-#include "iree/compiler/Dialect/HAL/Target/LLVM/LLVMIRPasses.h"
-#include "iree/compiler/Dialect/HAL/Target/TargetRegistry.h"
-#include "iree/compiler/Utils/FlatbufferUtils.h"
-#include "iree/schemas/llvmir_executable_def_builder.h"
-#include "llvm/IR/LLVMContext.h"
-#include "llvm/IR/Module.h"
-#include "llvm/Support/TargetSelect.h"
-#include "mlir/Target/LLVMIR.h"
-
-namespace mlir {
-namespace iree_compiler {
-namespace IREE {
-namespace HAL {
-
-class LLVMIRTargetBackend final : public LLVMBaseTargetBackend {
- public:
- explicit LLVMIRTargetBackend(LLVMTargetOptions options)
- : LLVMBaseTargetBackend(options) {}
-
- // NOTE: we could vary these based on the options, such as by arch/etc.
- std::string name() const override { return "llvm_ir"; }
- std::string filter_pattern() const override { return "llvm-ir*"; }
-
- LogicalResult serializeExecutable(IREE::HAL::ExecutableTargetOp targetOp,
- OpBuilder &executableBuilder) override {
- // Perform the translation to LLVM in a separate context to avoid
- // multi-threading issues.
- llvm::LLVMContext context;
- // Remove all private functions, e.g tile size calcuations.
- SmallVector<FuncOp, 4> nonPublicFn;
- for (auto func : targetOp.getInnerModule().getOps<FuncOp>()) {
- if (!func.isPublic()) {
- nonPublicFn.push_back(func);
- }
- }
- for (auto func : nonPublicFn) {
- func.erase();
- }
-
- // At this moment we are leaving MLIR LLVM dialect land translating module
- // into target independent LLVMIR.
- auto llvmModule =
- mlir::translateModuleToLLVMIR(targetOp.getInnerModule(), context);
-
- if (!llvmModule) {
- return targetOp.emitError("Failed to translate executable to LLVM IR");
- }
-
- // LLVMIR opt passes.
- auto targetMachine = createTargetMachine(options_);
- if (!targetMachine) {
- targetOp.emitError("Can't create target machine for target triple: " +
- options_.targetTriple);
- return failure();
- }
- LogicalResult translationResult =
- runLLVMIRPasses(options_, targetMachine.get(), llvmModule.get());
- if (failed(translationResult)) {
- return targetOp.emitError(
- "Can't build LLVMIR opt passes for ExecutableOp module");
- }
-
- // Serialize LLVM module directly into flatbuffer.
- FlatbufferBuilder builder;
- auto bitcodeModuleRef = builder.streamUint8Vec([&](raw_ostream &stream) {
- llvmModule->print(stream, nullptr);
- return true;
- });
-
- auto entryPointsRef = builder.createStringVec(llvm::map_range(
- targetOp.getBlock().getOps<ExecutableEntryPointOp>(),
- [&](ExecutableEntryPointOp op) { return op.sym_name(); }));
-
- iree_LLVMIRExecutableDef_start_as_root(builder);
- iree_LLVMIRExecutableDef_entry_points_add(builder, entryPointsRef);
- iree_LLVMIRExecutableDef_bitcode_module_add(builder, bitcodeModuleRef);
- iree_LLVMIRExecutableDef_end_as_root(builder);
-
- // Add the binary data to the target executable.
- executableBuilder.create<IREE::HAL::ExecutableBinaryOp>(
- targetOp.getLoc(),
- static_cast<uint32_t>(IREE::HAL::ExecutableFormat::LLVM),
- builder.getBufferAttr(executableBuilder.getContext()));
- return success();
- }
-};
-
-void registerLLVMIRTargetBackends(
- std::function<LLVMTargetOptions()> queryOptions) {
- getLLVMTargetOptionsFromFlags();
- static TargetBackendRegistration registration("llvm-ir", [=]() {
- // Initalize registered targets.
- llvm::InitializeNativeTarget();
- llvm::InitializeNativeTargetAsmPrinter();
- return std::make_unique<LLVMIRTargetBackend>(queryOptions());
- });
-}
-
-} // namespace HAL
-} // namespace IREE
-} // namespace iree_compiler
-} // namespace mlir
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/IR/LLVMIRTarget.h b/iree/compiler/Dialect/HAL/Target/LLVM/IR/LLVMIRTarget.h
deleted file mode 100644
index 46db6a9..0000000
--- a/iree/compiler/Dialect/HAL/Target/LLVM/IR/LLVMIRTarget.h
+++ /dev/null
@@ -1,35 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-//
-
-#ifndef IREE_COMPILER_DIALECT_HAL_TARGET_LLVM_IR_LLVMIRTARGET_H_
-#define IREE_COMPILER_DIALECT_HAL_TARGET_LLVM_IR_LLVMIRTARGET_H_
-
-#include "iree/compiler/Dialect/HAL/Target/LLVM/LLVMTargetOptions.h"
-
-namespace mlir {
-namespace iree_compiler {
-namespace IREE {
-namespace HAL {
-
-// Registers the LLVM IR target backends.
-void registerLLVMIRTargetBackends(
- std::function<LLVMTargetOptions()> queryOptions);
-
-} // namespace HAL
-} // namespace IREE
-} // namespace iree_compiler
-} // namespace mlir
-
-#endif // IREE_COMPILER_DIALECT_HAL_TARGET_LLVM_IR_LLVMIRTARGET_H_
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/LLVMBaseTarget.h b/iree/compiler/Dialect/HAL/Target/LLVM/LLVMBaseTarget.h
index 58dd82a..16b0a38 100644
--- a/iree/compiler/Dialect/HAL/Target/LLVM/LLVMBaseTarget.h
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/LLVMBaseTarget.h
@@ -23,7 +23,8 @@
namespace IREE {
namespace HAL {
-// Base target for LLVM ahead-of-time (AOT) and just-in-time (JIT) backends.
+// Base target for LLVM backends.
+// TODO(scotttodd): fold into LLVMAOTTarget now that LLVMJITTarget is gone
class LLVMBaseTargetBackend : public TargetBackend {
public:
explicit LLVMBaseTargetBackend(LLVMTargetOptions options);
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/LLVMTargetOptions.h b/iree/compiler/Dialect/HAL/Target/LLVM/LLVMTargetOptions.h
index ba454b6..1ac53bc 100644
--- a/iree/compiler/Dialect/HAL/Target/LLVM/LLVMTargetOptions.h
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/LLVMTargetOptions.h
@@ -48,8 +48,7 @@
bool keepLinkerArtifacts = false;
};
-// Returns LLVMTargetOptions struct intialized with the
-// iree-hal-llvm-ir-* flags.
+// Returns LLVMTargetOptions struct intialized with the iree-llvm-* flags.
LLVMTargetOptions getLLVMTargetOptionsFromFlags();
} // namespace HAL
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/test/binary_op.mlir b/iree/compiler/Dialect/HAL/Target/LLVM/test/binary_op.mlir
index 41f1939..ed951d7 100644
--- a/iree/compiler/Dialect/HAL/Target/LLVM/test/binary_op.mlir
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/test/binary_op.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -split-input-file -iree-hal-transformation-pipeline -iree-hal-target-backends=llvm-ir %s | IreeFileCheck %s
+// RUN: iree-opt -split-input-file -iree-hal-transformation-pipeline -iree-hal-target-backends=dylib-llvm-aot %s | IreeFileCheck %s
flow.executable @simpleMath_ex_dispatch_0 {
flow.dispatch.entry @simpleMath_rgn_dispatch_0 attributes {
workload = 4 : index
@@ -11,7 +11,7 @@
}
}
-// CHECK-LABEL: hal.executable @binary_op_linked_llvm_ir
+// CHECK-LABEL: hal.executable @binary_op_linked_llvm_aot
// CHECK-DAG: hal.executable.binary attributes {
// CHECK-SAME: data = dense
-// CHECK-SAME: format = 1280071245 : i32} {
+// CHECK-SAME: format = 1145850178 : i32} {
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/test/matmul_op.mlir b/iree/compiler/Dialect/HAL/Target/LLVM/test/matmul_op.mlir
index 1a48518..8d00275 100644
--- a/iree/compiler/Dialect/HAL/Target/LLVM/test/matmul_op.mlir
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/test/matmul_op.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -split-input-file -iree-hal-transformation-pipeline -iree-hal-target-backends=llvm-ir %s | IreeFileCheck %s
+// RUN: iree-opt -split-input-file -iree-hal-transformation-pipeline -iree-hal-target-backends=dylib-llvm-aot %s | IreeFileCheck %s
flow.executable @simpleMath_ex_dispatch_0 {
flow.dispatch.entry @simpleMath_rgn_dispatch_0 attributes {
workload = 4 : index
@@ -11,7 +11,7 @@
}
}
-// CHECK-LABEL: hal.executable @matmul_op_linked_llvm_ir
+// CHECK-LABEL: hal.executable @matmul_op_linked_llvm_aot
// CHECK-DAG: hal.executable.binary attributes {
// CHECK-SAME: data = dense
-// CHECK-SAME: format = 1280071245 : i32} {
+// CHECK-SAME: format = 1145850178 : i32} {
diff --git a/iree/hal/cts/semaphore_test.cc b/iree/hal/cts/semaphore_test.cc
index a26a5f1..385590c 100644
--- a/iree/hal/cts/semaphore_test.cc
+++ b/iree/hal/cts/semaphore_test.cc
@@ -110,8 +110,8 @@
// Waiting any semaphore to signal.
TEST_P(SemaphoreTest, WaitAny) {
// TODO: fix this.
- if (driver_->name() == "dylib" || driver_->name() == "llvmjit" ||
- driver_->name() == "vmla" || driver_->name() == "vulkan") {
+ if (driver_->name() == "dylib" || driver_->name() == "vmla" ||
+ driver_->name() == "vulkan") {
GTEST_SKIP();
}
diff --git a/iree/hal/drivers/BUILD b/iree/hal/drivers/BUILD
index c5eab06..8c421d7 100644
--- a/iree/hal/drivers/BUILD
+++ b/iree/hal/drivers/BUILD
@@ -28,7 +28,6 @@
] + [
# TODO(*): select() and only pull in based on build configuration.
"//iree/hal/dylib/registration",
- "//iree/hal/llvmjit/registration",
"//iree/hal/vmla/registration",
"//iree/hal/vulkan/registration",
],
diff --git a/iree/hal/drivers/CMakeLists.txt b/iree/hal/drivers/CMakeLists.txt
index 388f199..24961d6 100644
--- a/iree/hal/drivers/CMakeLists.txt
+++ b/iree/hal/drivers/CMakeLists.txt
@@ -18,9 +18,6 @@
if(${IREE_HAL_DRIVER_DYLIB})
list(APPEND IREE_HAL_DRIVER_MODULES iree::hal::dylib::registration)
endif()
-if(${IREE_HAL_DRIVER_LLVM})
- list(APPEND IREE_HAL_DRIVER_MODULES iree::hal::llvmjit::registration)
-endif()
if(${IREE_HAL_DRIVER_METAL})
list(APPEND IREE_HAL_DRIVER_MODULES iree::hal::metal::registration)
endif()
diff --git a/iree/hal/drivers/init.c b/iree/hal/drivers/init.c
index bae7588..2b9a2b9 100644
--- a/iree/hal/drivers/init.c
+++ b/iree/hal/drivers/init.c
@@ -20,10 +20,6 @@
#include "iree/hal/dylib/registration/driver_module.h"
#endif // IREE_HAL_HAVE_DYLIB_DRIVER_MODULE
-#if defined(IREE_HAL_HAVE_LLVMJIT_DRIVER_MODULE)
-#include "iree/hal/llvmjit/registration/driver_module.h"
-#endif // IREE_HAL_HAVE_LLVMJIT_DRIVER_MODULE
-
#if defined(IREE_HAL_HAVE_METAL_DRIVER_MODULE)
#include "iree/hal/metal/registration/driver_module.h"
#endif // IREE_HAL_HAVE_METAL_DRIVER_MODULE
@@ -45,11 +41,6 @@
z0, iree_hal_dylib_driver_module_register(registry));
#endif // IREE_HAL_HAVE_DYLIB_DRIVER_MODULE
-#if defined(IREE_HAL_HAVE_LLVMJIT_DRIVER_MODULE)
- IREE_RETURN_AND_END_ZONE_IF_ERROR(
- z0, iree_hal_llvmjit_driver_module_register(registry));
-#endif // IREE_HAL_HAVE_LLVMJIT_DRIVER_MODULE
-
#if defined(IREE_HAL_HAVE_METAL_DRIVER_MODULE)
IREE_RETURN_AND_END_ZONE_IF_ERROR(
z0, iree_hal_metal_driver_module_register(registry));
diff --git a/iree/hal/executable_format.h b/iree/hal/executable_format.h
index 03ba1f6..9664065 100644
--- a/iree/hal/executable_format.h
+++ b/iree/hal/executable_format.h
@@ -67,12 +67,6 @@
constexpr ExecutableFormat kExecutableFormatMetal =
MakeExecutableFormatID("MTLE");
-// LLVMIR executable in FlatBuffer format using the
-// https://github.com/google/iree/tree/main/iree/schemas/llvmir_executable_def.fbs
-// schema.
-constexpr ExecutableFormat kExecutableFormatLLVM =
- MakeExecutableFormatID("LLVM");
-
// Dynamic Library (dylib) executable in FlatBuffer format using the
// https://github.com/google/iree/tree/main/iree/schemas/dylib_executable_def.fbs
// schema
diff --git a/iree/hal/llvmjit/BUILD b/iree/hal/llvmjit/BUILD
deleted file mode 100644
index d7f5535..0000000
--- a/iree/hal/llvmjit/BUILD
+++ /dev/null
@@ -1,63 +0,0 @@
-# Copyright 2020 Google LLC
-#
-# Licensed under the Apache License, Version 2.0 (the "License");
-# you may not use this file except in compliance with the License.
-# You may obtain a copy of the License at
-#
-# https://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing, software
-# distributed under the License is distributed on an "AS IS" BASIS,
-# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-# See the License for the specific language governing permissions and
-# limitations under the License.
-
-# HAL implementation for jitting CPU code from LLVMIR.
-
-load("//iree:build_defs.oss.bzl", "iree_cmake_extra_content")
-
-package(
- default_visibility = ["//visibility:public"],
- features = ["layering_check"],
- licenses = ["notice"], # Apache 2.0
-)
-
-iree_cmake_extra_content(
- content = """
-if(NOT ${IREE_HAL_DRIVER_LLVM})
- return()
-endif()
-""",
-)
-
-cc_library(
- name = "llvmjit",
- srcs = [
- "llvmjit_device.cc",
- "llvmjit_driver.cc",
- "llvmjit_executable.cc",
- "llvmjit_executable_cache.cc",
- ],
- hdrs = [
- "llvmjit_device.h",
- "llvmjit_driver.h",
- "llvmjit_executable.h",
- "llvmjit_executable_cache.h",
- ],
- deps = [
- "//iree/base:flatcc",
- "//iree/base:status",
- "//iree/base:tracing",
- "//iree/hal",
- "//iree/hal/host:host_executable",
- "//iree/hal/host:host_local_device",
- "//iree/hal/host/serial:serial_scheduling_model",
- "//iree/schemas:llvmir_executable_def_c_fbs",
- "@com_google_absl//absl/types:span",
- "@llvm-project//llvm:AsmParser",
- "@llvm-project//llvm:Core",
- "@llvm-project//llvm:ExecutionEngine",
- "@llvm-project//llvm:OrcJIT",
- "@llvm-project//llvm:Support",
- ],
-)
diff --git a/iree/hal/llvmjit/CMakeLists.txt b/iree/hal/llvmjit/CMakeLists.txt
deleted file mode 100644
index ea303d1..0000000
--- a/iree/hal/llvmjit/CMakeLists.txt
+++ /dev/null
@@ -1,50 +0,0 @@
-# Copyright 2020 Google LLC
-#
-# Licensed under the Apache License, Version 2.0 (the "License");
-# you may not use this file except in compliance with the License.
-# You may obtain a copy of the License at
-#
-# https://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing, software
-# distributed under the License is distributed on an "AS IS" BASIS,
-# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-# See the License for the specific language governing permissions and
-# limitations under the License.
-
-if(NOT ${IREE_HAL_DRIVER_LLVM})
- return()
-endif()
-
-iree_add_all_subdirs()
-
-iree_cc_library(
- NAME
- llvmjit
- HDRS
- "llvmjit_device.h"
- "llvmjit_driver.h"
- "llvmjit_executable.h"
- "llvmjit_executable_cache.h"
- SRCS
- "llvmjit_device.cc"
- "llvmjit_driver.cc"
- "llvmjit_executable.cc"
- "llvmjit_executable_cache.cc"
- DEPS
- LLVMAsmParser
- LLVMCore
- LLVMExecutionEngine
- LLVMOrcJIT
- LLVMSupport
- absl::span
- iree::base::flatcc
- iree::base::status
- iree::base::tracing
- iree::hal
- iree::hal::host::host_executable
- iree::hal::host::host_local_device
- iree::hal::host::serial::serial_scheduling_model
- iree::schemas::llvmir_executable_def_c_fbs
- PUBLIC
-)
diff --git a/iree/hal/llvmjit/llvmjit_device.cc b/iree/hal/llvmjit/llvmjit_device.cc
deleted file mode 100644
index 1a4c801..0000000
--- a/iree/hal/llvmjit/llvmjit_device.cc
+++ /dev/null
@@ -1,40 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#include "iree/hal/llvmjit/llvmjit_device.h"
-
-#include <utility>
-
-#include "iree/base/tracing.h"
-#include "iree/hal/llvmjit/llvmjit_executable_cache.h"
-
-namespace iree {
-namespace hal {
-namespace llvmjit {
-
-LLVMJITDevice::LLVMJITDevice(
- DeviceInfo device_info,
- std::unique_ptr<host::SchedulingModel> scheduling_model)
- : HostLocalDevice(std::move(device_info), std::move(scheduling_model)) {}
-
-LLVMJITDevice::~LLVMJITDevice() = default;
-
-ref_ptr<ExecutableCache> LLVMJITDevice::CreateExecutableCache() {
- IREE_TRACE_SCOPE0("LLVMJITDevice::CreateExecutableCache");
- return make_ref<LLVMJITExecutableCache>();
-}
-
-} // namespace llvmjit
-} // namespace hal
-} // namespace iree
diff --git a/iree/hal/llvmjit/llvmjit_device.h b/iree/hal/llvmjit/llvmjit_device.h
deleted file mode 100644
index d2f48d5..0000000
--- a/iree/hal/llvmjit/llvmjit_device.h
+++ /dev/null
@@ -1,37 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#ifndef IREE_HAL_LLVMJIT_LLVMJIT_DEVICE_H_
-#define IREE_HAL_LLVMJIT_LLVMJIT_DEVICE_H_
-
-#include "iree/hal/host/host_local_device.h"
-
-namespace iree {
-namespace hal {
-namespace llvmjit {
-
-class LLVMJITDevice final : public host::HostLocalDevice {
- public:
- LLVMJITDevice(DeviceInfo device_info,
- std::unique_ptr<host::SchedulingModel> scheduling_model);
- ~LLVMJITDevice() override;
-
- ref_ptr<ExecutableCache> CreateExecutableCache() override;
-};
-
-} // namespace llvmjit
-} // namespace hal
-} // namespace iree
-
-#endif // IREE_HAL_LLVMJIT_LLVMJIT_DEVICE_H_
diff --git a/iree/hal/llvmjit/llvmjit_driver.cc b/iree/hal/llvmjit/llvmjit_driver.cc
deleted file mode 100644
index 81077a6..0000000
--- a/iree/hal/llvmjit/llvmjit_driver.cc
+++ /dev/null
@@ -1,64 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#include "iree/hal/llvmjit/llvmjit_driver.h"
-
-#include <memory>
-
-#include "iree/hal/device_info.h"
-#include "iree/hal/host/serial/serial_scheduling_model.h"
-#include "iree/hal/llvmjit/llvmjit_device.h"
-
-namespace iree {
-namespace hal {
-namespace llvmjit {
-namespace {
-
-DeviceInfo GetDefaultDeviceInfo() {
- DeviceFeatureBitfield supported_features = DeviceFeature::kNone;
- // TODO(benvanik): implement debugging/profiling features.
- // supported_features |= DeviceFeature::kDebugging;
- // supported_features |= DeviceFeature::kCoverage;
- // supported_features |= DeviceFeature::kProfiling;
- DeviceInfo device_info("llvm-ir-jit", "llvm", supported_features);
- // TODO(benvanik): device info.
- return device_info;
-}
-
-} // namespace
-
-LLVMJITDriver::LLVMJITDriver() : Driver("llvmjit") {}
-
-LLVMJITDriver::~LLVMJITDriver() = default;
-
-StatusOr<std::vector<DeviceInfo>> LLVMJITDriver::EnumerateAvailableDevices() {
- std::vector<DeviceInfo> device_infos;
- device_infos.push_back(GetDefaultDeviceInfo());
- return device_infos;
-}
-
-StatusOr<ref_ptr<Device>> LLVMJITDriver::CreateDefaultDevice() {
- return CreateDevice(0);
-}
-
-StatusOr<ref_ptr<Device>> LLVMJITDriver::CreateDevice(
- DriverDeviceID device_id) {
- auto scheduling_model = std::make_unique<host::SerialSchedulingModel>();
- return make_ref<LLVMJITDevice>(GetDefaultDeviceInfo(),
- std::move(scheduling_model));
-}
-
-} // namespace llvmjit
-} // namespace hal
-} // namespace iree
diff --git a/iree/hal/llvmjit/llvmjit_driver.h b/iree/hal/llvmjit/llvmjit_driver.h
deleted file mode 100644
index fb903c8..0000000
--- a/iree/hal/llvmjit/llvmjit_driver.h
+++ /dev/null
@@ -1,40 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#ifndef IREE_HAL_LLVMJIT_LLVMJIT_DRIVER_H_
-#define IREE_HAL_LLVMJIT_LLVMJIT_DRIVER_H_
-
-#include "iree/hal/driver.h"
-
-namespace iree {
-namespace hal {
-namespace llvmjit {
-
-class LLVMJITDriver final : public Driver {
- public:
- LLVMJITDriver();
- ~LLVMJITDriver() override;
-
- StatusOr<std::vector<DeviceInfo>> EnumerateAvailableDevices() override;
-
- StatusOr<ref_ptr<Device>> CreateDefaultDevice() override;
-
- StatusOr<ref_ptr<Device>> CreateDevice(DriverDeviceID device_id) override;
-};
-
-} // namespace llvmjit
-} // namespace hal
-} // namespace iree
-
-#endif // IREE_HAL_LLVMJIT_LLVMJIT_DRIVER_H_
diff --git a/iree/hal/llvmjit/llvmjit_executable.cc b/iree/hal/llvmjit/llvmjit_executable.cc
deleted file mode 100644
index 2db4bd8..0000000
--- a/iree/hal/llvmjit/llvmjit_executable.cc
+++ /dev/null
@@ -1,236 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#include "iree/hal/llvmjit/llvmjit_executable.h"
-
-#include <iostream>
-#include <memory>
-
-#include "iree/base/tracing.h"
-#include "iree/hal/buffer.h"
-#include "iree/hal/executable.h"
-#include "llvm/ADT/ArrayRef.h"
-#include "llvm/ADT/StringRef.h"
-#include "llvm/AsmParser/Parser.h"
-#include "llvm/ExecutionEngine/Orc/ThreadSafeModule.h"
-#include "llvm/IR/LLVMContext.h"
-#include "llvm/Support/Error.h"
-#include "llvm/Support/MemoryBuffer.h"
-#include "llvm/Support/SourceMgr.h"
-
-// flatcc schemas:
-#include "iree/base/flatcc.h"
-#include "iree/schemas/llvmir_executable_def_reader.h"
-#include "iree/schemas/llvmir_executable_def_verifier.h"
-
-// NOTE: starting to port this to C.
-
-// Verifies the structure of the flatbuffer so that we can avoid doing so during
-// runtime. There are still some conditions we must be aware of (such as omitted
-// names on functions with internal linkage), however we shouldn't need to
-// bounds check anything within the flatbuffer after this succeeds.
-static iree_status_t iree_hal_llvmir_executable_flatbuffer_verify(
- iree_const_byte_span_t flatbuffer_data) {
- if (!flatbuffer_data.data || flatbuffer_data.data_length < 16) {
- return iree_make_status(
- IREE_STATUS_INVALID_ARGUMENT,
- "flatbuffer data is not present or less than 16 bytes (%zu total)",
- flatbuffer_data.data_length);
- }
-
- // Run flatcc generated verification. This ensures all pointers are in-bounds
- // and that we can safely walk the file, but not that the actual contents of
- // the flatbuffer meet our expectations.
- int verify_ret = iree_LLVMIRExecutableDef_verify_as_root(
- flatbuffer_data.data, flatbuffer_data.data_length);
- if (verify_ret != flatcc_verify_ok) {
- return iree_make_status(IREE_STATUS_INVALID_ARGUMENT,
- "flatbuffer verification failed: %s",
- flatcc_verify_error_string(verify_ret));
- }
-
- iree_LLVMIRExecutableDef_table_t executable_def =
- iree_LLVMIRExecutableDef_as_root(flatbuffer_data.data);
-
- flatbuffers_string_vec_t entry_points_vec =
- iree_LLVMIRExecutableDef_entry_points_get(executable_def);
- size_t entry_point_count = flatbuffers_string_vec_len(entry_points_vec);
- for (size_t i = 0; i < entry_point_count; ++i) {
- if (!flatbuffers_string_len(
- flatbuffers_string_vec_at(entry_points_vec, i))) {
- return iree_make_status(IREE_STATUS_INVALID_ARGUMENT,
- "executable entry point %zu has no name", i);
- }
- }
-
- if (!flatbuffers_uint8_vec_len(
- iree_LLVMIRExecutableDef_bitcode_module_get(executable_def))) {
- return iree_make_status(IREE_STATUS_INVALID_ARGUMENT,
- "executable bitcode_module is missing/empty");
- }
-
- return iree_ok_status();
-}
-
-namespace iree {
-namespace hal {
-namespace llvmjit {
-
-// static
-StatusOr<ref_ptr<LLVMJITExecutable>> LLVMJITExecutable::Load(
- ExecutableSpec spec, bool allow_aliasing_data) {
- IREE_TRACE_SCOPE0("LLVMJITExecutable::Load");
-
- // Verify and fetch the executable flatbuffer wrapper.
- iree_const_byte_span_t executable_data = iree_make_const_byte_span(
- spec.executable_data.data(), spec.executable_data.size());
- IREE_RETURN_IF_ERROR(
- iree_hal_llvmir_executable_flatbuffer_verify(executable_data));
- iree_LLVMIRExecutableDef_table_t executable_def =
- iree_LLVMIRExecutableDef_as_root(executable_data.data);
-
- flatbuffers_uint8_vec_t bitcode_module_vec =
- iree_LLVMIRExecutableDef_bitcode_module_get(executable_def);
- auto mem_buffer = llvm::MemoryBuffer::getMemBufferCopy(
- llvm::StringRef(reinterpret_cast<const char*>(bitcode_module_vec),
- flatbuffers_uint8_vec_len(bitcode_module_vec)),
- "llvm-ir");
- auto llvm_context = std::make_unique<llvm::LLVMContext>();
- llvm::SMDiagnostic sm_diagnostic;
- auto module = llvm::parseAssembly(*mem_buffer, sm_diagnostic, *llvm_context);
- if (!module) {
- return InvalidArgumentErrorBuilder(IREE_LOC)
- << "Can't parse LLVMIR Module: " << sm_diagnostic.getMessage().str();
- }
- auto dataLayout = module->getDataLayout();
- llvm::orc::ThreadSafeModule thread_safe_module(std::move(module),
- std::move(llvm_context));
- auto ll_jit = llvm::cantFail(llvm::orc::LLJITBuilder().create());
-
- llvm::Error err = ll_jit->addIRModule(std::move(thread_safe_module));
- if (err) {
- return InvalidArgumentErrorBuilder(IREE_LOC)
- << "Can't add executable module to executable LLJIT"
- << llvm::toString(std::move(err));
- }
-
- auto llvmjit_serarch_generator =
- llvm::orc::DynamicLibrarySearchGenerator::GetForCurrentProcess(
- dataLayout.getGlobalPrefix());
- if (!llvmjit_serarch_generator) {
- return UnavailableErrorBuilder(IREE_LOC)
- << "Can't resolve symbols in current process: "
- << llvm::toString(llvmjit_serarch_generator.takeError());
- }
-
- auto& main_jitllvmjit = ll_jit->getMainJITDylib();
- main_jitllvmjit.addGenerator(std::move(llvmjit_serarch_generator.get()));
-
- auto executable =
- make_ref<LLVMJITExecutable>(spec, std::move(ll_jit), allow_aliasing_data);
-
- flatbuffers_string_vec_t entry_points =
- iree_LLVMIRExecutableDef_entry_points_get(executable_def);
- executable->symbols_.resize(flatbuffers_string_vec_len(entry_points));
- for (size_t i = 0; i < flatbuffers_string_vec_len(entry_points); ++i) {
- flatbuffers_string_t entry_point =
- flatbuffers_string_vec_at(entry_points, i);
- auto func_symbol = executable->ll_jit_->lookup(
- llvm::StringRef(entry_point, flatbuffers_string_len(entry_point)));
- if (!func_symbol) {
- return NotFoundErrorBuilder(IREE_LOC)
- << "Can't JIT compile function '" << entry_point
- << "': " << llvm::toString(func_symbol.takeError());
- }
- executable->symbols_[i] = func_symbol.get();
- }
-
- return executable;
-}
-
-LLVMJITExecutable::LLVMJITExecutable(ExecutableSpec spec,
- std::unique_ptr<llvm::orc::LLJIT> ll_jit,
- bool allow_aliasing_data)
- : spec_(spec), ll_jit_(std::move(ll_jit)) {
- if (!allow_aliasing_data) {
- // Clone data.
- cloned_executable_data_ = {spec.executable_data.begin(),
- spec.executable_data.end()};
- spec_.executable_data = absl::MakeConstSpan(cloned_executable_data_);
- }
-}
-
-LLVMJITExecutable::~LLVMJITExecutable() = default;
-
-struct LLVMJITDispatchState : public HostExecutable::DispatchState {
- LLVMJITDispatchState() = default;
-
- std::array<uint32_t, 3> workgroup_count;
- std::array<uint32_t, 3> workgroup_size;
- llvm::JITEvaluatedSymbol symbol;
- llvm::SmallVector<void*, 4> args;
- llvm::SmallVector<int32_t, 4> push_constant;
-};
-
-StatusOr<ref_ptr<HostExecutable::DispatchState>>
-LLVMJITExecutable::PrepareDispatch(const DispatchParams& params) {
- IREE_TRACE_SCOPE0("LLVMJITExecutable::PrepareDispatch");
-
- if (params.entry_point >= symbols_.size()) {
- return InvalidArgumentErrorBuilder(IREE_LOC)
- << "Invalid entry point ordinal " << params.entry_point;
- }
-
- auto dispatch_state = make_ref<LLVMJITDispatchState>();
- dispatch_state->workgroup_count = params.workgroup_count;
- dispatch_state->workgroup_size = params.workgroup_size;
- dispatch_state->symbol = symbols_[params.entry_point];
-
- for (size_t set = 0; set < params.set_bindings.size(); ++set) {
- for (size_t binding = 0; binding < params.set_bindings[set].size();
- ++binding) {
- const auto& io_binding = params.set_bindings[set][binding];
- IREE_ASSIGN_OR_RETURN(auto memory,
- io_binding.buffer->MapMemory<uint8_t>(
- MemoryAccessBitfield::kWrite, io_binding.offset,
- io_binding.length));
- auto data = memory.mutable_data();
- dispatch_state->args.push_back(data);
- }
- }
- // TODO(ataei): Consider moving this casting to codegen side ?!
- for (int i = 0; i < params.push_constants->values.size(); ++i) {
- dispatch_state->push_constant.push_back(params.push_constants->values[i]);
- }
-
- return std::move(dispatch_state);
-}
-
-Status LLVMJITExecutable::DispatchTile(DispatchState* state,
- std::array<uint32_t, 3> workgroup_xyz) {
- IREE_TRACE_SCOPE0("LLVMJITExecutable::DispatchTile");
- auto* dispatch_state = static_cast<LLVMJITDispatchState*>(state);
-
- auto func_ptr = (void (*)(void**, int32_t*, uint32_t*, uint32_t*,
- uint32_t*))dispatch_state->symbol.getAddress();
- func_ptr(dispatch_state->args.data(), dispatch_state->push_constant.data(),
- workgroup_xyz.data(), dispatch_state->workgroup_count.data(),
- dispatch_state->workgroup_size.data());
-
- return OkStatus();
-}
-
-} // namespace llvmjit
-} // namespace hal
-} // namespace iree
diff --git a/iree/hal/llvmjit/llvmjit_executable.h b/iree/hal/llvmjit/llvmjit_executable.h
deleted file mode 100644
index 6c5b3bc..0000000
--- a/iree/hal/llvmjit/llvmjit_executable.h
+++ /dev/null
@@ -1,60 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#ifndef IREE_HAL_LLVMJIT_LLVMJIT_EXECUTABLE_H_
-#define IREE_HAL_LLVMJIT_LLVMJIT_EXECUTABLE_H_
-
-#include <vector>
-
-#include "iree/base/status.h"
-#include "iree/hal/executable_spec.h"
-#include "iree/hal/host/host_executable.h"
-#include "llvm/ADT/ArrayRef.h"
-#include "llvm/ExecutionEngine/Orc/LLJIT.h"
-
-namespace iree {
-namespace hal {
-namespace llvmjit {
-
-struct MemrefType;
-
-class LLVMJITExecutable final : public HostExecutable {
- public:
- static StatusOr<ref_ptr<LLVMJITExecutable>> Load(ExecutableSpec spec,
- bool allow_aliasing_data);
-
- LLVMJITExecutable(ExecutableSpec spec,
- std::unique_ptr<llvm::orc::LLJIT> ll_jit,
- bool allow_aliasing_data);
- ~LLVMJITExecutable() override;
-
- bool supports_debugging() const override { return false; }
-
- StatusOr<ref_ptr<DispatchState>> PrepareDispatch(
- const DispatchParams& params) override;
- Status DispatchTile(DispatchState* state,
- std::array<uint32_t, 3> workgroup_xyz) override;
-
- private:
- ExecutableSpec spec_;
- std::vector<uint8_t> cloned_executable_data_;
- std::unique_ptr<llvm::orc::LLJIT> ll_jit_;
- llvm::SmallVector<llvm::JITEvaluatedSymbol, 4> symbols_;
-};
-
-} // namespace llvmjit
-} // namespace hal
-} // namespace iree
-
-#endif // IREE_HAL_LLVMJIT_LLVMJIT_EXECUTABLE_H_
diff --git a/iree/hal/llvmjit/llvmjit_executable_cache.cc b/iree/hal/llvmjit/llvmjit_executable_cache.cc
deleted file mode 100644
index 401cd90..0000000
--- a/iree/hal/llvmjit/llvmjit_executable_cache.cc
+++ /dev/null
@@ -1,50 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#include "iree/hal/llvmjit/llvmjit_executable_cache.h"
-
-#include "iree/base/status.h"
-#include "iree/base/tracing.h"
-#include "iree/hal/executable_format.h"
-#include "iree/hal/llvmjit/llvmjit_executable.h"
-
-namespace iree {
-namespace hal {
-namespace llvmjit {
-
-LLVMJITExecutableCache::LLVMJITExecutableCache() = default;
-
-LLVMJITExecutableCache::~LLVMJITExecutableCache() = default;
-
-bool LLVMJITExecutableCache::CanPrepareFormat(ExecutableFormat format) const {
- return format == kExecutableFormatLLVM;
-}
-
-StatusOr<ref_ptr<Executable>> LLVMJITExecutableCache::PrepareExecutable(
- ExecutableLayout* executable_layout, ExecutableCachingModeBitfield mode,
- const ExecutableSpec& spec) {
- IREE_TRACE_SCOPE0("LLVMJITExecutableCache::PrepareExecutable");
-
- // Wrap the data (or copy it).
- bool allow_aliasing_data =
- AllBitsSet(mode, ExecutableCachingMode::kAliasProvidedData);
- IREE_ASSIGN_OR_RETURN(auto executable,
- LLVMJITExecutable::Load(spec, !allow_aliasing_data));
-
- return executable;
-}
-
-} // namespace llvmjit
-} // namespace hal
-} // namespace iree
diff --git a/iree/hal/llvmjit/llvmjit_executable_cache.h b/iree/hal/llvmjit/llvmjit_executable_cache.h
deleted file mode 100644
index 0725423..0000000
--- a/iree/hal/llvmjit/llvmjit_executable_cache.h
+++ /dev/null
@@ -1,41 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#ifndef IREE_HAL_LLVMJIT_EXECUTABLE_CACHE_H_
-#define IREE_HAL_LLVMJIT_EXECUTABLE_CACHE_H_
-
-#include "iree/hal/executable.h"
-#include "iree/hal/executable_cache.h"
-
-namespace iree {
-namespace hal {
-namespace llvmjit {
-
-class LLVMJITExecutableCache final : public ExecutableCache {
- public:
- LLVMJITExecutableCache();
- ~LLVMJITExecutableCache() override;
-
- bool CanPrepareFormat(ExecutableFormat format) const override;
-
- StatusOr<ref_ptr<Executable>> PrepareExecutable(
- ExecutableLayout* executable_layout, ExecutableCachingModeBitfield mode,
- const ExecutableSpec& spec) override;
-};
-
-} // namespace llvmjit
-} // namespace hal
-} // namespace iree
-
-#endif // IREE_HAL_LLVMJIT_EXECUTABLE_CACHE_H_
diff --git a/iree/hal/llvmjit/registration/BUILD b/iree/hal/llvmjit/registration/BUILD
deleted file mode 100644
index 5d37108..0000000
--- a/iree/hal/llvmjit/registration/BUILD
+++ /dev/null
@@ -1,53 +0,0 @@
-# Copyright 2020 Google LLC
-#
-# Licensed under the Apache License, Version 2.0 (the "License");
-# you may not use this file except in compliance with the License.
-# You may obtain a copy of the License at
-#
-# https://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing, software
-# distributed under the License is distributed on an "AS IS" BASIS,
-# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-# See the License for the specific language governing permissions and
-# limitations under the License.
-
-load("//iree:build_defs.oss.bzl", "iree_cmake_extra_content")
-
-package(
- default_visibility = ["//visibility:public"],
- features = ["layering_check"],
- licenses = ["notice"], # Apache 2.0
-)
-
-iree_cmake_extra_content(
- content = """
-if(${IREE_HAL_DRIVER_LLVM})
-""",
- inline = True,
-)
-
-cc_library(
- name = "registration",
- srcs = ["driver_module.cc"],
- hdrs = ["driver_module.h"],
- defines = [
- "IREE_HAL_HAVE_LLVMJIT_DRIVER_MODULE=1",
- ],
- deps = [
- "//iree/base:flags",
- "//iree/base:status",
- "//iree/hal:api",
- "//iree/hal/llvmjit",
- "@llvm-project//llvm:Support",
- # TODO(ataei): Link with native target dep.
- "@llvm-project//llvm:X86CodeGen",
- ],
-)
-
-iree_cmake_extra_content(
- content = """
-endif()
-""",
- inline = True,
-)
diff --git a/iree/hal/llvmjit/registration/CMakeLists.txt b/iree/hal/llvmjit/registration/CMakeLists.txt
deleted file mode 100644
index cab37cb..0000000
--- a/iree/hal/llvmjit/registration/CMakeLists.txt
+++ /dev/null
@@ -1,38 +0,0 @@
-# Copyright 2020 Google LLC
-#
-# Licensed under the Apache License, Version 2.0 (the "License");
-# you may not use this file except in compliance with the License.
-# You may obtain a copy of the License at
-#
-# https://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing, software
-# distributed under the License is distributed on an "AS IS" BASIS,
-# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-# See the License for the specific language governing permissions and
-# limitations under the License.
-
-iree_add_all_subdirs()
-
-if(${IREE_HAL_DRIVER_LLVM})
-
-iree_cc_library(
- NAME
- registration
- HDRS
- "driver_module.h"
- SRCS
- "driver_module.cc"
- DEPS
- LLVMSupport
- LLVMX86CodeGen
- iree::base::flags
- iree::base::status
- iree::hal::api
- iree::hal::llvmjit
- DEFINES
- "IREE_HAL_HAVE_LLVMJIT_DRIVER_MODULE=1"
- PUBLIC
-)
-
-endif()
diff --git a/iree/hal/llvmjit/registration/driver_module.cc b/iree/hal/llvmjit/registration/driver_module.cc
deleted file mode 100644
index 65d2047..0000000
--- a/iree/hal/llvmjit/registration/driver_module.cc
+++ /dev/null
@@ -1,61 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#include "iree/hal/llvmjit/registration/driver_module.h"
-
-#include <inttypes.h>
-
-#include "iree/hal/llvmjit/llvmjit_driver.h"
-#include "llvm/Support/TargetSelect.h"
-
-#define IREE_HAL_LLVMJIT_DRIVER_ID 0x4C4C564Du // LLVM
-
-static iree_status_t iree_hal_llvmjit_driver_factory_enumerate(
- void* self, const iree_hal_driver_info_t** out_driver_infos,
- iree_host_size_t* out_driver_info_count) {
- static const iree_hal_driver_info_t driver_infos[1] = {{
- /*driver_id=*/IREE_HAL_LLVMJIT_DRIVER_ID,
- /*driver_name=*/iree_make_cstring_view("llvm"),
- /*full_name=*/iree_make_cstring_view("LLVM Bitcode JIT (deprecated)"),
- }};
- *out_driver_info_count = IREE_ARRAYSIZE(driver_infos);
- *out_driver_infos = driver_infos;
- return iree_ok_status();
-}
-
-static iree_status_t iree_hal_llvmjit_driver_factory_try_create(
- void* self, iree_hal_driver_id_t driver_id, iree_allocator_t allocator,
- iree_hal_driver_t** out_driver) {
- if (driver_id != IREE_HAL_LLVMJIT_DRIVER_ID) {
- return iree_make_status(IREE_STATUS_UNAVAILABLE,
- "no driver with ID %016" PRIu64
- " is provided by this factory",
- driver_id);
- }
- llvm::InitializeNativeTarget();
- llvm::InitializeNativeTargetAsmPrinter();
- auto* driver = new iree::hal::llvmjit::LLVMJITDriver();
- *out_driver = reinterpret_cast<iree_hal_driver_t*>(driver);
- return iree_ok_status();
-}
-
-IREE_API_EXPORT iree_status_t IREE_API_CALL
-iree_hal_llvmjit_driver_module_register(iree_hal_driver_registry_t* registry) {
- static const iree_hal_driver_factory_t factory = {
- /*self=*/NULL,
- iree_hal_llvmjit_driver_factory_enumerate,
- iree_hal_llvmjit_driver_factory_try_create,
- };
- return iree_hal_driver_registry_register_factory(registry, &factory);
-}
diff --git a/iree/hal/llvmjit/registration/driver_module.h b/iree/hal/llvmjit/registration/driver_module.h
deleted file mode 100644
index 0300782..0000000
--- a/iree/hal/llvmjit/registration/driver_module.h
+++ /dev/null
@@ -1,31 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#ifndef IREE_HAL_LLVMJIT_REGISTRATION_DRIVER_MODULE_H_
-#define IREE_HAL_LLVMJIT_REGISTRATION_DRIVER_MODULE_H_
-
-#include "iree/hal/api.h"
-
-#ifdef __cplusplus
-extern "C" {
-#endif // __cplusplus
-
-IREE_API_EXPORT iree_status_t IREE_API_CALL
-iree_hal_llvmjit_driver_module_register(iree_hal_driver_registry_t* registry);
-
-#ifdef __cplusplus
-} // extern "C"
-#endif // __cplusplus
-
-#endif // IREE_HAL_LLVMJIT_REGISTRATION_DRIVER_MODULE_H_
diff --git a/iree/schemas/BUILD b/iree/schemas/BUILD
index 98fd357..2d9676f 100644
--- a/iree/schemas/BUILD
+++ b/iree/schemas/BUILD
@@ -41,12 +41,6 @@
)
iree_flatbuffer_c_library(
- name = "llvmir_executable_def_c_fbs",
- srcs = ["llvmir_executable_def.fbs"],
- flatcc_args = FLATCC_ARGS,
-)
-
-iree_flatbuffer_c_library(
name = "metal_executable_def_c_fbs",
srcs = ["metal_executable_def.fbs"],
flatcc_args = FLATCC_ARGS,
@@ -69,7 +63,6 @@
targets = [
":bytecode_module_def_c_fbs",
":dylib_executable_def_c_fbs",
- ":llvmir_executable_def_c_fbs",
":metal_executable_def_c_fbs",
":spirv_executable_def_c_fbs",
":vmla_executable_def_c_fbs",
diff --git a/iree/schemas/CMakeLists.txt b/iree/schemas/CMakeLists.txt
index 68d435c..b49b3bb 100644
--- a/iree/schemas/CMakeLists.txt
+++ b/iree/schemas/CMakeLists.txt
@@ -42,19 +42,6 @@
flatbuffer_c_library(
NAME
- llvmir_executable_def_c_fbs
- SRCS
- "llvmir_executable_def.fbs"
- FLATCC_ARGS
- "--reader"
- "--builder"
- "--verifier"
- "--json"
- PUBLIC
-)
-
-flatbuffer_c_library(
- NAME
metal_executable_def_c_fbs
SRCS
"metal_executable_def.fbs"
diff --git a/iree/schemas/llvmir_executable_def.fbs b/iree/schemas/llvmir_executable_def.fbs
deleted file mode 100644
index 5897900..0000000
--- a/iree/schemas/llvmir_executable_def.fbs
+++ /dev/null
@@ -1,31 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-namespace iree;
-
-// 'LLVMIR Executable'.
-
-file_identifier "LLVM";
-file_extension "ll";
-
-// Machine independent LLVMIR executable module.
-// This exeuctable will be compiled with the target machine later on.
-table LLVMIRExecutableDef {
- // A map of entry points to string names with the same order as in the executable op.
- entry_points:[string];
- // A serialized llvm::Module object.
- bitcode_module:[ubyte];
-}
-
-root_type LLVMIRExecutableDef;
diff --git a/iree/test/e2e/models/edge_detection.mlir b/iree/test/e2e/models/edge_detection.mlir
index c0cfcbe..b8f7b73 100644
--- a/iree/test/e2e/models/edge_detection.mlir
+++ b/iree/test/e2e/models/edge_detection.mlir
@@ -1,7 +1,7 @@
// RUN: iree-run-mlir -export-all -iree-hal-target-backends=vmla %s -function-input="1x128x128x1xf32" | IreeFileCheck %s
// RUN: [[ $IREE_LLVMAOT_DISABLE == 1 ]] || (iree-run-mlir -export-all -iree-hal-target-backends=dylib-llvm-aot %s -function-input="1x128x128x1xf32" | IreeFileCheck %s)
// RUN: [[ $IREE_VULKAN_DISABLE == 1 ]] || (iree-run-mlir -export-all -iree-hal-target-backends=vulkan-spirv %s -function-input="1x128x128x1xf32" | IreeFileCheck %s)
-// RUN: [[ $IREE_LLVMJIT_DISABLE == 1 ]] || (iree-run-mlir -export-all -iree-hal-target-backends=llvm-ir -iree-enable-linalg-on-tensors %s -function-input="1x128x128x1xf32" | IreeFileCheck %s)
+// RUN: [[ $IREE_LLVMAOT_DISABLE == 1 ]] || (iree-run-mlir -export-all -iree-hal-target-backends=dylib-llvm-aot -iree-enable-linalg-on-tensors %s -function-input="1x128x128x1xf32" | IreeFileCheck %s)
// RUN: [[ $IREE_VULKAN_DISABLE == 1 ]] || (iree-run-mlir -export-all -iree-hal-target-backends=vulkan-spirv -iree-enable-linalg-on-tensors %s -function-input="1x128x128x1xf32" | IreeFileCheck %s)
// Image edge detection module generated by iree/colab/edge_detection.ipynb.
diff --git a/iree/test/e2e/models/fragment_000.mlir b/iree/test/e2e/models/fragment_000.mlir
index 2786c6b..4e6b895 100644
--- a/iree/test/e2e/models/fragment_000.mlir
+++ b/iree/test/e2e/models/fragment_000.mlir
@@ -1,7 +1,7 @@
// RUN: iree-run-mlir -export-all -iree-hal-target-backends=vmla %s -function-input="f32=0" -function-input="5x1xf32=[1][-2][-3][4][-5]" -function-input="f32=1" -function-input="5x5xf32=[3.46499 -7.64389 -5.72249 5.98053 17.6892][2.9707 -6.20734 -4.25962 4.76055 13.8784][2.47641 -4.77079 -2.79675 3.54056 10.0675][1.98212 -3.33424 -1.33388 2.32058 6.25666][1.48783 -1.8977 0.12899 1.1006 2.4458]" -function-input="5xf32=0 0 0 0 0" | IreeFileCheck %s --implicit-check-not="[" --implicit-check-not="]"
// RUN: [[ $IREE_LLVMAOT_DISABLE == 1 ]] || (iree-run-mlir -export-all -iree-hal-target-backends=dylib-llvm-aot %s -function-input="f32=0" -function-input="5x1xf32=[1][-2][-3][4][-5]" -function-input="f32=1" -function-input="5x5xf32=[3.46499 -7.64389 -5.72249 5.98053 17.6892][2.9707 -6.20734 -4.25962 4.76055 13.8784][2.47641 -4.77079 -2.79675 3.54056 10.0675][1.98212 -3.33424 -1.33388 2.32058 6.25666][1.48783 -1.8977 0.12899 1.1006 2.4458]" -function-input="5xf32=0 0 0 0 0" | IreeFileCheck %s --implicit-check-not="[" --implicit-check-not="]")
// RUN: [[ $IREE_VULKAN_DISABLE == 1 ]] || (iree-run-mlir -export-all -iree-hal-target-backends=vulkan-spirv %s -function-input="f32=0" -function-input="5x1xf32=[1][-2][-3][4][-5]" -function-input="f32=1" -function-input="5x5xf32=[3.46499 -7.64389 -5.72249 5.98053 17.6892][2.9707 -6.20734 -4.25962 4.76055 13.8784][2.47641 -4.77079 -2.79675 3.54056 10.0675][1.98212 -3.33424 -1.33388 2.32058 6.25666][1.48783 -1.8977 0.12899 1.1006 2.4458]" -function-input="5xf32=0 0 0 0 0" | IreeFileCheck %s --implicit-check-not="[" --implicit-check-not="]")
-// RUN: [[ $IREE_LLVMJIT_DISABLE == 1 ]] || (iree-run-mlir -export-all -iree-hal-target-backends=llvm-ir -iree-enable-linalg-on-tensors %s -function-input="f32=0" -function-input="5x1xf32=[1][-2][-3][4][-5]" -function-input="f32=1" -function-input="5x5xf32=[3.46499 -7.64389 -5.72249 5.98053 17.6892][2.9707 -6.20734 -4.25962 4.76055 13.8784][2.47641 -4.77079 -2.79675 3.54056 10.0675][1.98212 -3.33424 -1.33388 2.32058 6.25666][1.48783 -1.8977 0.12899 1.1006 2.4458]" -function-input="5xf32=0 0 0 0 0" | IreeFileCheck %s --implicit-check-not="[" --implicit-check-not="]")
+// RUN: [[ $IREE_LLVMAOT_DISABLE == 1 ]] || (iree-run-mlir -export-all -iree-hal-target-backends=dylib-llvm-aot -iree-enable-linalg-on-tensors %s -function-input="f32=0" -function-input="5x1xf32=[1][-2][-3][4][-5]" -function-input="f32=1" -function-input="5x5xf32=[3.46499 -7.64389 -5.72249 5.98053 17.6892][2.9707 -6.20734 -4.25962 4.76055 13.8784][2.47641 -4.77079 -2.79675 3.54056 10.0675][1.98212 -3.33424 -1.33388 2.32058 6.25666][1.48783 -1.8977 0.12899 1.1006 2.4458]" -function-input="5xf32=0 0 0 0 0" | IreeFileCheck %s --implicit-check-not="[" --implicit-check-not="]")
// RUN: [[ $IREE_VULKAN_DISABLE == 1 ]] || (iree-run-mlir -export-all -iree-hal-target-backends=vulkan-spirv -iree-enable-linalg-on-tensors %s -function-input="f32=0" -function-input="5x1xf32=[1][-2][-3][4][-5]" -function-input="f32=1" -function-input="5x5xf32=[3.46499 -7.64389 -5.72249 5.98053 17.6892][2.9707 -6.20734 -4.25962 4.76055 13.8784][2.47641 -4.77079 -2.79675 3.54056 10.0675][1.98212 -3.33424 -1.33388 2.32058 6.25666][1.48783 -1.8977 0.12899 1.1006 2.4458]" -function-input="5xf32=0 0 0 0 0" | IreeFileCheck %s --implicit-check-not="[" --implicit-check-not="]")
// CHECK-LABEL: EXEC @main_entry_dispatch_3
diff --git a/iree/test/e2e/models/fullyconnected.mlir b/iree/test/e2e/models/fullyconnected.mlir
index a7ccbfd..e99f401 100644
--- a/iree/test/e2e/models/fullyconnected.mlir
+++ b/iree/test/e2e/models/fullyconnected.mlir
@@ -1,7 +1,7 @@
// RUN: iree-run-mlir -export-all %s -iree-hal-target-backends=vmla -function-input="1x5xf32=1,-2,-3,4,-5" -function-input="1x5x3x1xf32=15,14,13,12,11,10,9,8,7,6,5,4,3,2,1" | IreeFileCheck %s
// RUN: [[ $IREE_LLVMAOT_DISABLE == 1 ]] || (iree-run-mlir -export-all %s -iree-hal-target-backends=dylib-llvm-aot -function-input="1x5xf32=1,-2,-3,4,-5" -function-input="1x5x3x1xf32=15,14,13,12,11,10,9,8,7,6,5,4,3,2,1" | IreeFileCheck %s)
// RUN: [[ $IREE_VULKAN_DISABLE == 1 ]] || (iree-run-mlir -export-all %s -iree-hal-target-backends=vulkan-spirv -function-input="1x5xf32=1,-2,-3,4,-5" -function-input="1x5x3x1xf32=15,14,13,12,11,10,9,8,7,6,5,4,3,2,1" | IreeFileCheck %s)
-// RUN: [[ $IREE_LLVMJIT_DISABLE == 1 ]] || (iree-run-mlir -export-all %s -iree-hal-target-backends=llvm-ir -function-input="1x5xf32=1,-2,-3,4,-5" -function-input="1x5x3x1xf32=15,14,13,12,11,10,9,8,7,6,5,4,3,2,1" -iree-enable-linalg-on-tensors | IreeFileCheck %s)
+// RUN: [[ $IREE_LLVMAOT_DISABLE == 1 ]] || (iree-run-mlir -export-all %s -iree-hal-target-backends=dylib-llvm-aot -function-input="1x5xf32=1,-2,-3,4,-5" -function-input="1x5x3x1xf32=15,14,13,12,11,10,9,8,7,6,5,4,3,2,1" -iree-enable-linalg-on-tensors | IreeFileCheck %s)
// RUN: [[ $IREE_VULKAN_DISABLE == 1 ]] || (iree-run-mlir -export-all %s -iree-hal-target-backends=vulkan-spirv -function-input="1x5xf32=1,-2,-3,4,-5" -function-input="1x5x3x1xf32=15,14,13,12,11,10,9,8,7,6,5,4,3,2,1" -iree-enable-linalg-on-tensors | IreeFileCheck %s)
// CHECK-LABEL: EXEC @main
diff --git a/iree/test/e2e/models/mnist_fake_weights.mlir b/iree/test/e2e/models/mnist_fake_weights.mlir
index 29e8bf2..caafffe 100644
--- a/iree/test/e2e/models/mnist_fake_weights.mlir
+++ b/iree/test/e2e/models/mnist_fake_weights.mlir
@@ -3,7 +3,7 @@
// RUN: iree-run-mlir -export-all -iree-hal-target-backends=vmla %s -function-input="1x28x28x1xf32" | IreeFileCheck %s
// RUN: [[ $IREE_LLVMAOT_DISABLE == 1 ]] || (iree-run-mlir -export-all -iree-hal-target-backends=dylib-llvm-aot %s -function-input="1x28x28x1xf32" | IreeFileCheck %s)
// RUN: [[ $IREE_VULKAN_DISABLE == 1 ]] || (iree-run-mlir -export-all -iree-hal-target-backends=vulkan-spirv %s -function-input="1x28x28x1xf32" | IreeFileCheck %s)
-// RUN: [[ $IREE_LLVMJIT_DISABLE == 1 ]] || (iree-run-mlir -export-all -iree-hal-target-backends=llvm-ir -iree-enable-linalg-on-tensors %s -function-input="1x28x28x1xf32" | IreeFileCheck %s)
+// RUN: [[ $IREE_LLVMAOT_DISABLE == 1 ]] || (iree-run-mlir -export-all -iree-hal-target-backends=dylib-llvm-aot -iree-enable-linalg-on-tensors %s -function-input="1x28x28x1xf32" | IreeFileCheck %s)
// RUN: [[ $IREE_VULKAN_DISABLE == 1 ]] || (iree-run-mlir -export-all -iree-hal-target-backends=vulkan-spirv -iree-enable-linalg-on-tensors %s -function-input="1x28x28x1xf32" | IreeFileCheck %s)
module {
diff --git a/iree/test/e2e/models/unidirectional_lstm.mlir b/iree/test/e2e/models/unidirectional_lstm.mlir
index 0dc9929..926078d 100644
--- a/iree/test/e2e/models/unidirectional_lstm.mlir
+++ b/iree/test/e2e/models/unidirectional_lstm.mlir
@@ -3,7 +3,7 @@
// RUN: iree-run-mlir %s -iree-hal-target-backends=vmla -function-input="1x5xf32=[0 1 0 3 4]" -function-input="1x5x2x2xf32=[1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20]" | IreeFileCheck %s --implicit-check-not="[" --implicit-check-not="]"
// RUN: [[ $IREE_LLVMAOT_DISABLE == 1 ]] || (iree-run-mlir %s -iree-hal-target-backends=dylib-llvm-aot -function-input="1x5xf32=[0 1 0 3 4]" -function-input="1x5x2x2xf32=[1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20]" | IreeFileCheck %s --implicit-check-not="[" --implicit-check-not="]")
// RUN: [[ $IREE_VULKAN_DISABLE == 1 ]] || (iree-run-mlir %s -iree-hal-target-backends=vulkan-spirv -function-input="1x5xf32=[0 1 0 3 4]" -function-input="1x5x2x2xf32=[1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20]" | IreeFileCheck %s --implicit-check-not="[" --implicit-check-not="]")
-// RUN: [[ $IREE_LLVMJIT_DISABLE == 1 ]] || (iree-run-mlir %s -iree-hal-target-backends=llvm-ir -function-input="1x5xf32=[0 1 0 3 4]" -function-input="1x5x2x2xf32=[1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20]" -iree-enable-linalg-on-tensors | IreeFileCheck %s --implicit-check-not="[" --implicit-check-not="]")
+// RUN: [[ $IREE_LLVMAOT_DISABLE == 1 ]] || (iree-run-mlir %s -iree-hal-target-backends=dylib-llvm-aot -function-input="1x5xf32=[0 1 0 3 4]" -function-input="1x5x2x2xf32=[1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20]" -iree-enable-linalg-on-tensors | IreeFileCheck %s --implicit-check-not="[" --implicit-check-not="]")
// RUN: [[ $IREE_VULKAN_DISABLE == 1 ]] || (iree-run-mlir %s -iree-hal-target-backends=vulkan-spirv -function-input="1x5xf32=[0 1 0 3 4]" -function-input="1x5x2x2xf32=[1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20]" -iree-enable-linalg-on-tensors | IreeFileCheck %s --implicit-check-not="[" --implicit-check-not="]")
// Exported via the XLA HLO Importer
diff --git a/iree/tools/BUILD b/iree/tools/BUILD
index fb581b1..309d028 100644
--- a/iree/tools/BUILD
+++ b/iree/tools/BUILD
@@ -154,13 +154,11 @@
hdrs = ["init_targets.h"],
local_defines = [
"IREE_HAVE_LLVMAOT_TARGET",
- "IREE_HAVE_LLVMIR_TARGET",
"IREE_HAVE_VMLA_TARGET",
"IREE_HAVE_VULKANSPIRV_TARGET",
],
deps = [
"//iree/compiler/Dialect/HAL/Target/LLVM/AOT:LLVMAOT",
- "//iree/compiler/Dialect/HAL/Target/LLVM/IR:LLVMIR",
"//iree/compiler/Dialect/HAL/Target/VMLA",
"//iree/compiler/Dialect/HAL/Target/VulkanSPIRV",
],
diff --git a/iree/tools/CMakeLists.txt b/iree/tools/CMakeLists.txt
index f4a8e02..3d7b447 100644
--- a/iree/tools/CMakeLists.txt
+++ b/iree/tools/CMakeLists.txt
@@ -25,10 +25,6 @@
list(APPEND IREE_COMPILER_TARGETS iree::compiler::Dialect::HAL::Target::LLVM::AOT::LLVMAOT)
list(APPEND IREE_COMPILER_TARGET_COPTS "-DIREE_HAVE_LLVMAOT_TARGET")
endif()
-if(${IREE_TARGET_BACKEND_LLVM-IR})
- list(APPEND IREE_COMPILER_TARGETS iree::compiler::Dialect::HAL::Target::LLVM::IR::LLVMIR)
- list(APPEND IREE_COMPILER_TARGET_COPTS "-DIREE_HAVE_LLVMIR_TARGET")
-endif()
if(${IREE_TARGET_BACKEND_METAL-SPIRV})
list(APPEND IREE_COMPILER_TARGETS iree::compiler::Dialect::HAL::Target::MetalSPIRV)
list(APPEND IREE_COMPILER_TARGET_COPTS "-DIREE_HAVE_METALSPIRV_TARGET")
diff --git a/iree/tools/init_targets.cc b/iree/tools/init_targets.cc
index 509461d..e16548e 100644
--- a/iree/tools/init_targets.cc
+++ b/iree/tools/init_targets.cc
@@ -17,9 +17,6 @@
#ifdef IREE_HAVE_LLVMAOT_TARGET
#include "iree/compiler/Dialect/HAL/Target/LLVM/AOT/LLVMAOTTarget.h"
#endif
-#ifdef IREE_HAVE_LLVMIR_TARGET
-#include "iree/compiler/Dialect/HAL/Target/LLVM/IR/LLVMIRTarget.h"
-#endif
#ifdef IREE_HAVE_VMLA_TARGET
#include "iree/compiler/Dialect/HAL/Target/VMLA/VMLATarget.h"
#endif
@@ -44,10 +41,6 @@
IREE::HAL::registerLLVMAOTTargetBackends(
[]() { return IREE::HAL::getLLVMTargetOptionsFromFlags(); });
#endif
-#ifdef IREE_HAVE_LLVMIR_TARGET
- IREE::HAL::registerLLVMIRTargetBackends(
- []() { return IREE::HAL::getLLVMTargetOptionsFromFlags(); });
-#endif
#ifdef IREE_HAVE_VMLA_TARGET
IREE::HAL::registerVMLATargetBackends(
[]() { return IREE::HAL::getVMLATargetOptionsFromFlags(); });
