Merge pull request #3901 from GMNGeoffrey:main-to-google
PiperOrigin-RevId: 343112174
diff --git a/docs/get_started/getting_started_android_cmake.md b/docs/get_started/getting_started_android_cmake.md
index 05cbd71..874b183 100644
--- a/docs/get_started/getting_started_android_cmake.md
+++ b/docs/get_started/getting_started_android_cmake.md
@@ -276,8 +276,8 @@
```shell
$ ../iree-build-android/host/bin/iree-translate \
-iree-mlir-to-vm-bytecode-module \
- -iree-llvm-target-triple=aarch64-linux-android \
-iree-hal-target-backends=dylib-llvm-aot \
+ -iree-llvm-target-triple=aarch64-linux-android \
$PWD/iree/tools/test/simple.mlir \
-o /tmp/simple-llvm_aot.vmfb
```
diff --git a/docs/get_started/getting_started_windows_cmake.md b/docs/get_started/getting_started_windows_cmake.md
index a6025b2..04a48e0 100644
--- a/docs/get_started/getting_started_windows_cmake.md
+++ b/docs/get_started/getting_started_windows_cmake.md
@@ -89,6 +89,51 @@
> cmake --build ..\iree-build\
```
+## Target Configuration
+
+### LLVM AOT Backend
+
+`-iree-hal-target-backends=dylib-llvm-aot` can be used to generate modules with
+ahead-of-time compiled kernels stored in DLLs. Run the iree-opt/iree-translate
+tools from a command prompt with `lld-link.exe` or `link.exe` tools on the
+`PATH` and the MSVC/Windows SDK environment variables; the easiest way to get
+this configured is to use the `vsvarsall.bat` or `vcvars64.bat` files to set
+your environment. See
+[the Microsoft documentation](https://docs.microsoft.com/en-us/cpp/build/building-on-the-command-line?view=vs-2019)
+for details on configuring the toolchain.
+
+If you want to manually specify the linker used, set the
+`IREE_LLVMAOT_LINKER_PATH` environment variable to the path of the linker:
+
+```powershell
+> set IREE_LLVMAOT_LINKER_PATH="C:\Tools\LLVM\bin\lld-link.exe"
+```
+
+Translate a source MLIR file into an IREE module:
+
+```powershell
+> ..\iree-build\iree\tools\iree-translate.exe \
+ -iree-mlir-to-vm-bytecode-module \
+ -iree-hal-target-backends=dylib-llvm-aot \
+ iree/tools/test/simple.mlir \
+ -o %TMP%/simple-llvm_aot.vmfb
+```
+
+Note that this will use the host machine as the target by default, and the
+exact target triple and architecture can be specified with flags when
+cross-compiling:
+
+```powershell
+> ..\iree-build\iree\tools\iree-translate.exe \
+ -iree-mlir-to-vm-bytecode-module \
+ -iree-hal-target-backends=dylib-llvm-aot \
+ -iree-llvm-target-triple=x86_64-pc-windows-msvc \
+ -iree-llvm-target-cpu=host \
+ -iree-llvm-target-cpu-features=host \
+ iree/tools/test/simple.mlir \
+ -o %TMP%/simple-llvm_aot.vmfb
+```
+
## What's next?
### Take a Look Around
diff --git a/integrations/tensorflow/bindings/python/pyiree/tf/support/BUILD b/integrations/tensorflow/bindings/python/pyiree/tf/support/BUILD
index fc67c96..9b1c869 100644
--- a/integrations/tensorflow/bindings/python/pyiree/tf/support/BUILD
+++ b/integrations/tensorflow/bindings/python/pyiree/tf/support/BUILD
@@ -60,7 +60,6 @@
],
python_version = "PY3",
tags = [
- "driver=llvm",
"driver=vmla",
],
deps = INTREE_TENSORFLOW_PY_DEPS + [
diff --git a/integrations/tensorflow/bindings/python/pyiree/tf/support/tf_utils.py b/integrations/tensorflow/bindings/python/pyiree/tf/support/tf_utils.py
index 187b765..591df43 100644
--- a/integrations/tensorflow/bindings/python/pyiree/tf/support/tf_utils.py
+++ b/integrations/tensorflow/bindings/python/pyiree/tf/support/tf_utils.py
@@ -992,11 +992,6 @@
"driver": "vmla",
"compiler_targets": ["vmla"]
},
- "iree_llvmjit": {
- "compiled_module_class": IreeCompiledModule,
- "driver": "llvm",
- "compiler_targets": ["llvm-ir"]
- },
"iree_vulkan": {
"compiled_module_class": IreeCompiledModule,
"driver": "vulkan",
@@ -1009,13 +1004,13 @@
Args:
backend_name: a str specifying which backend to use. Should be one of
- 'tf', 'iree_vmla', 'iree_llvmjit', 'iree_vulkan'.
+ 'tf', 'iree_vmla', 'iree_vulkan'.
backend_id: an optional str specifying what name to use when saving
compiled artifacts. Must satisfy `backend_id.startswith(backend_name)`.
Raises:
- KeyError: if backend_name is not one of ['tf', 'iree_vmla',
- 'iree_llvmjit', 'iree_vulkan'].
+ KeyError: if backend_name is not one of
+ ['tf', 'iree_vmla', 'iree_vulkan'].
ValueError: if backend_id doesn't start with backend_name.
"""
if backend_name not in self._name_to_info:
diff --git a/integrations/tensorflow/e2e/README.md b/integrations/tensorflow/e2e/README.md
index 0ac1605..4b457b5 100644
--- a/integrations/tensorflow/e2e/README.md
+++ b/integrations/tensorflow/e2e/README.md
@@ -18,7 +18,7 @@
If you do not have your environment setup to use IREE with Vulkan (see
[this doc](https://google.github.io/iree/get-started/generic-vulkan-env-setup)),
then you can run the manual test targets with
-`--target_backends=tf,iree_vmla,iree_llvmjit` (that is, by omitting
+`--target_backends=tf,iree_vmla` (that is, by omitting
`iree_vulkan` from the list of backends to run the tests on).
The test suites can be run excluding Vulkan by specifying
@@ -234,8 +234,8 @@
to check numerical correctness against TensorFlow. Tests targets that pass are
placed into the `e2e_tests` test suite. Tests that fail on particular backends
are recorded in lists in the `BUILD` files. For example, if
-`experimental_new_test.py` fails on the `iree_llvmjit` and `iree_vulkan`
-backends then the following lines should be added to the `BUILD` file:
+`experimental_new_test.py` fails on the `iree_vulkan` backend then the following
+lines should be added to the `BUILD` file:
```build
LLVM_FAILING = [
diff --git a/integrations/tensorflow/e2e/iree_e2e_cartesian_product_test_suite.bzl b/integrations/tensorflow/e2e/iree_e2e_cartesian_product_test_suite.bzl
index ee6d4f1..7f49f83 100644
--- a/integrations/tensorflow/e2e/iree_e2e_cartesian_product_test_suite.bzl
+++ b/integrations/tensorflow/e2e/iree_e2e_cartesian_product_test_suite.bzl
@@ -164,6 +164,13 @@
tests = []
for flags in all_flag_configurations:
+ if len(flags["target_backends"].split(",")) > 1:
+ fail("Multiple target backends cannot be specified at once, but " +
+ "got `{}`".format(flags["target_backends"]))
+ driver = get_driver(flags["target_backends"])
+ if not driver:
+ continue
+
# Check if this is a failing configuration.
failing = flags in failing_flag_configurations
@@ -180,12 +187,6 @@
tests.append(test_name)
args = ["--{}={}".format(k, v) for k, v in flags.items()]
-
- if len(flags["target_backends"].split(",")) > 1:
- fail("Multiple target backends cannot be specified at once, but " +
- "got `{}`".format(flags["target_backends"]))
-
- driver = get_driver(flags["target_backends"])
py_test_tags = ["driver={}".format(driver)]
if tags != None: # `is` is not supported.
py_test_tags += tags
diff --git a/integrations/tensorflow/e2e/iree_e2e_test_suite.bzl b/integrations/tensorflow/e2e/iree_e2e_test_suite.bzl
index 8fb35a4..58be0f6 100644
--- a/integrations/tensorflow/e2e/iree_e2e_test_suite.bzl
+++ b/integrations/tensorflow/e2e/iree_e2e_test_suite.bzl
@@ -19,8 +19,10 @@
def get_driver(backend):
# TODO(#2175): Simplify this after backend names are standardized.
driver = backend.replace("iree_", "") # "iree_<driver>" --> "<driver>"
+
+ # TODO(#2673): enable LLVM AOT for these tests. JIT is deprecated.
if driver == "llvmjit":
- driver = "llvm"
+ driver = ""
return driver
def iree_e2e_test_suite(
@@ -70,6 +72,8 @@
]
driver = get_driver(backend)
+ if not driver:
+ continue
py_test_tags = ["driver={}".format(driver)]
if tags != None: # `is` is not supported.
py_test_tags += tags
diff --git a/integrations/tensorflow/e2e/keras/BUILD b/integrations/tensorflow/e2e/keras/BUILD
index d6e61ad..468eb79 100644
--- a/integrations/tensorflow/e2e/keras/BUILD
+++ b/integrations/tensorflow/e2e/keras/BUILD
@@ -44,13 +44,13 @@
vision_model_test_manual is for manual testing of all keras vision models.
Test will run only manually with all parameters specified manually, for example:
bazel run -c opt integrations/tensorflow/e2e/keras:vision_model_test_manual -- \
---target_backends=tf,iree_vmla,iree_llvmjit \
+--target_backends=tf,iree_vmla \
--data=imagenet \
--url=https://storage.googleapis.com/iree_models/ \
--model=ResNet50
Command arguments description:
---target_backends: can be combination of these: tf,iree_vmla,iree_llvmjit
+--target_backends: can be combination of these: tf,iree_vmla
--data: can be 'imagenet' or 'cifar10'.
imagenet - input image size (1, 224, 224, 3)
cifar10 - input image size (1, 32, 32, 3) - it is used for quick tests
diff --git a/integrations/tensorflow/e2e/slim_vision_models/BUILD b/integrations/tensorflow/e2e/slim_vision_models/BUILD
index 88f60b0..881aff3 100644
--- a/integrations/tensorflow/e2e/slim_vision_models/BUILD
+++ b/integrations/tensorflow/e2e/slim_vision_models/BUILD
@@ -76,7 +76,6 @@
"resnet_v2_152",
],
"target_backends": [
- "iree_llvmjit",
"iree_vulkan",
],
},
@@ -159,7 +158,6 @@
"tf",
"tflite",
"iree_vmla",
- "iree_llvmjit",
"iree_vulkan",
],
},
diff --git a/iree/base/wait_handle.h b/iree/base/wait_handle.h
index adff1c9..b95bf7e 100644
--- a/iree/base/wait_handle.h
+++ b/iree/base/wait_handle.h
@@ -32,8 +32,7 @@
#endif // IREE_PLATFORM_*
// TODO(benvanik): see if we can get sync file on linux too:
-// TODO(scotttodd): fix include on Android (missing linkopts?)
-#if 0 && defined(IREE_PLATFORM_ANDROID)
+#if defined(IREE_PLATFORM_ANDROID)
#define IREE_HAVE_WAIT_TYPE_SYNC_FILE 1
#endif // IREE_PLATFORM_ANDROID
diff --git a/iree/base/wait_handle_posix.c b/iree/base/wait_handle_posix.c
index 8f180d1..088f6cc 100644
--- a/iree/base/wait_handle_posix.c
+++ b/iree/base/wait_handle_posix.c
@@ -27,7 +27,7 @@
#include <sys/eventfd.h>
#endif // IREE_HAVE_WAIT_TYPE_EVENTFD
#if defined(IREE_HAVE_WAIT_TYPE_SYNC_FILE)
-#include <sync.h>
+#include <android/sync.h>
#endif // IREE_HAVE_WAIT_TYPE_SYNC_FILE
//===----------------------------------------------------------------------===//
diff --git a/iree/compiler/Conversion/CodegenUtils/GetNumWorkgroups.h b/iree/compiler/Conversion/CodegenUtils/GetNumWorkgroups.h
index e8cd9db..01e4b3d 100644
--- a/iree/compiler/Conversion/CodegenUtils/GetNumWorkgroups.h
+++ b/iree/compiler/Conversion/CodegenUtils/GetNumWorkgroups.h
@@ -28,7 +28,7 @@
namespace mlir {
class Location;
class FuncOp;
-class LogicalResult;
+struct LogicalResult;
class PatternRewriter;
class ConversionPatternRewriter;
class Value;
diff --git a/iree/compiler/Conversion/HLOToLinalg/HLOToLinalgOnBuffers.cpp b/iree/compiler/Conversion/HLOToLinalg/HLOToLinalgOnBuffers.cpp
index 9f33bd1..7381818 100644
--- a/iree/compiler/Conversion/HLOToLinalg/HLOToLinalgOnBuffers.cpp
+++ b/iree/compiler/Conversion/HLOToLinalg/HLOToLinalgOnBuffers.cpp
@@ -665,44 +665,55 @@
namespace {
/// Converts mhlo.slice operation to linalg.subview + linalg.copy
-struct SliceOpConversion
- : public ConvertToLinalgBufferOp<SliceOpConversion, mhlo::SliceOp> {
- using ConvertToLinalgBufferOp<SliceOpConversion,
- mhlo::SliceOp>::ConvertToLinalgBufferOp;
+struct SliceOpConversion : public OpConversionPattern<mhlo::SliceOp> {
+ SliceOpConversion(MLIRContext *context,
+ TensorToBufferMap const &resultTensorToBufferMap,
+ PatternBenefit benefit = 1)
+ : OpConversionPattern<mhlo::SliceOp>(context, benefit),
+ resultTensorToBufferMap(resultTensorToBufferMap) {}
- LogicalResult apply(mhlo::SliceOp op, ArrayRef<Value> inputBuffers,
- ArrayRef<Value> resultBuffers,
- ConversionPatternRewriter &rewriter) const;
+ LogicalResult matchAndRewrite(
+ mhlo::SliceOp op, ArrayRef<Value> inputBuffers,
+ ConversionPatternRewriter &rewriter) const override {
+ auto loc = op.getLoc();
+ auto argType = inputBuffers[0].getType().template dyn_cast<ShapedType>();
+ if (!argType || !argType.hasStaticShape()) {
+ return op.emitError("expected static shape");
+ }
+
+ auto resultShape = op.getResult().getType().cast<ShapedType>().getShape();
+ SmallVector<Value, 3> offsets, sizes, strides;
+ for (int i = 0, e = argType.getRank(); i < e; ++i) {
+ Value startIndex = rewriter.create<ConstantIndexOp>(
+ loc, op.start_indices().getValue<int64_t>(i));
+ offsets.push_back(startIndex);
+ Value size = rewriter.create<ConstantIndexOp>(loc, resultShape[i]);
+ sizes.push_back(size);
+ Value stride = rewriter.create<ConstantIndexOp>(
+ loc, op.strides().getValue<int64_t>(i));
+ strides.push_back(stride);
+ }
+ auto subViewOp = rewriter.create<SubViewOp>(loc, inputBuffers[0], offsets,
+ sizes, strides);
+
+ // If the result of the subview is already mapped to a buffer, a copy is
+ // required from the buffer above into the mapped buffer.
+ if (Value bufferForResult =
+ resultTensorToBufferMap.lookup(op.getResult())) {
+ rewriter.create<linalg::CopyOp>(loc, subViewOp, bufferForResult);
+ rewriter.replaceOp(op, bufferForResult);
+ } else {
+ rewriter.replaceOp(op, subViewOp.getResult());
+ }
+
+ return success();
+ }
+
+ private:
+ TensorToBufferMap const &resultTensorToBufferMap;
};
} // namespace
-LogicalResult SliceOpConversion::apply(
- mhlo::SliceOp op, ArrayRef<Value> inputBuffers,
- ArrayRef<Value> resultBuffers, ConversionPatternRewriter &rewriter) const {
- auto loc = op.getLoc();
- auto argType = inputBuffers[0].getType().template dyn_cast<ShapedType>();
- if (!argType || !argType.hasRank()) {
- return op.emitError("expected known-rank args");
- }
-
- SmallVector<Value, 3> offsets, sizes, strides;
- for (int i = 0, e = argType.getRank(); i < e; ++i) {
- Value startIndex = rewriter.create<ConstantIndexOp>(
- loc, op.start_indices().getValue<int64_t>(i));
- offsets.push_back(startIndex);
- Value size = rewriter.create<DimOp>(loc, resultBuffers[0], i);
- sizes.push_back(size);
- Value stride = rewriter.create<ConstantIndexOp>(
- loc, op.strides().getValue<int64_t>(i));
- strides.push_back(stride);
- }
- auto subViewOp =
- rewriter.create<SubViewOp>(loc, inputBuffers[0], offsets, sizes, strides);
- rewriter.create<linalg::CopyOp>(loc, subViewOp, resultBuffers[0]);
-
- return success();
-}
-
//===----------------------------------------------------------------------===//
// mhlo.reduce_window conversion patterns and utility functions.
//===----------------------------------------------------------------------===//
diff --git a/iree/compiler/Conversion/HLOToLinalg/test/slice.mlir b/iree/compiler/Conversion/HLOToLinalg/test/slice.mlir
index 406a62b..0de43b3 100644
--- a/iree/compiler/Conversion/HLOToLinalg/test/slice.mlir
+++ b/iree/compiler/Conversion/HLOToLinalg/test/slice.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -split-input-file -iree-codegen-hlo-to-linalg-on-buffers -cse %s | IreeFileCheck %s
+// RUN: iree-opt -split-input-file -iree-codegen-hlo-to-linalg-pipeline -canonicalize %s | IreeFileCheck %s
module {
// CHECK_LABEL: @slice_whole_buffer
@@ -25,19 +25,10 @@
// -----
module {
- // CHECK: #[[MAP:.+]] = affine_map<(d0, d1)[s0, s1, s2] -> (d0 * s1 + s0 + d1 * s2)>
+ // CHECK: #[[MAP:.+]] = affine_map<(d0, d1) -> (d0 * 4 + d1 + 4)>
// CHECK: @slice_whole_stride
- // CHECK-DAG: %[[OUT:.+]] = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<1x4xi32>
// CHECK-DAG: %[[IN:.+]] = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<3x4xi32>
- // CHECK-DAG: %[[ZERO:.+]] = constant 0 : index
- // CHECK-DAG: %[[ONE:.+]] = constant 1 : index
- // CHECK-DAG: %[[DIM0:.+]] = dim %[[OUT]], %[[ZERO]] : memref<1x4xi32>
- // CHECK-DAG: %[[DIM1:.+]] = dim %[[OUT]], %[[ONE]] : memref<1x4xi32>
- // CHECK: subview %[[IN]]
- // CHECK-SAME: [%[[ONE]], %[[ZERO]]]
- // CHECK-SAME: [%[[DIM0]], %[[DIM1]]]
- // CHECK-SAME: [%[[ONE]], %[[ONE]]]
- // CHECK-SAME: : memref<3x4xi32> to memref<?x?xi32, #[[MAP]]>
+ // CHECK: subview %[[IN]][1, 0] [1, 4] [1, 1] : memref<3x4xi32> to memref<1x4xi32, #[[MAP]]>
// CHECK: linalg.copy
func @slice_whole_stride()
attributes {signature = (tensor<3x4xi32>) -> (tensor<1x4xi32>)} {
@@ -60,19 +51,10 @@
// -----
module {
- // CHECK: #[[MAP:.+]] = affine_map<(d0, d1)[s0, s1, s2] -> (d0 * s1 + s0 + d1 * s2)>
+ // CHECK: #[[MAP:.+]] = affine_map<(d0, d1) -> (d0 * 4 + d1 + 5)>
// CHECK: @slice_stride_part
- // CHECK-DAG: %[[OUT:.+]] = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<1x2xi32>
// CHECK-DAG: %[[IN:.+]] = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<3x4xi32>
- // CHECK-DAG: %[[ZERO:.+]] = constant 0 : index
- // CHECK-DAG: %[[ONE:.+]] = constant 1 : index
- // CHECK-DAG: %[[DIM0:.+]] = dim %[[OUT]], %[[ZERO]] : memref<1x2xi32>
- // CHECK-DAG: %[[DIM1:.+]] = dim %[[OUT]], %[[ONE]] : memref<1x2xi32>
- // CHECK: subview %[[IN]]
- // CHECK-SAME: [%[[ONE]], %[[ONE]]]
- // CHECK-SAME: [%[[DIM0]], %[[DIM1]]]
- // CHECK-SAME: [%[[ONE]], %[[ONE]]]
- // CHECK-SAME: : memref<3x4xi32> to memref<?x?xi32, #map>
+ // CHECK: subview %[[IN]][1, 1] [1, 2] [1, 1] : memref<3x4xi32> to memref<1x2xi32, #[[MAP]]>
// CHECK: linalg.copy
func @slice_stride_part()
attributes {signature = (tensor<3x4xi32>) -> (tensor<1x2xi32>)} {
@@ -91,3 +73,32 @@
hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write"
}
}
+
+// -----
+
+// CHECK-DAG: #[[MAP0:.+]] = affine_map<(d0, d1) -> (d0 * 4 + d1 + 5)>
+// CHECK-DAG: #[[MAP1:.+]] = affine_map<(d0, d1) -> (d0, d1)>
+// CHECK: func @slice_stride_part
+// CHECK: %[[OUT:.+]] = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0, operand_result_index = 2 : i32} : memref<1x2xi32>
+// CHECK: %[[IN0:.+]] = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0, operand_result_index = 0 : i32} : memref<3x4xi32>
+// CHECK: %[[SUBVIEW:.+]] = subview %[[IN0]][1, 1] [1, 2] [1, 1] : memref<3x4xi32> to memref<1x2xi32, #[[MAP0]]>
+// CHECK: %[[IN1:.+]] = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg1, operand_result_index = 1 : i32} : memref<1x2xi32>
+// CHECK: linalg.generic
+// CHECK-SAME: ins(%[[SUBVIEW]], %[[IN1]]
+// CHECK-SAME: outs(%[[OUT]]
+module {
+ func @slice_stride_part() {
+ %c0 = constant 0 : index
+ %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 {operand_result_index = 0 : i32} : tensor<3x4xi32>
+ %1 = hal.interface.load.tensor @legacy_io::@arg1, offset = %c0 {operand_result_index = 1 : i32} : tensor<1x2xi32>
+ %2 = "mhlo.slice"(%0) {limit_indices = dense<[2, 3]> : tensor<2xi64>, start_indices = dense<1> : tensor<2xi64>, strides = dense<1> : tensor<2xi64>} : (tensor<3x4xi32>) -> tensor<1x2xi32>
+ %3 = mhlo.add %2, %1 : tensor<1x2xi32>
+ hal.interface.store.tensor %3, @legacy_io::@ret0, offset = %c0 {operand_result_index = 2 : i32} : tensor<1x2xi32>
+ return
+ }
+ hal.interface @legacy_io attributes {sym_visibility = "private"} {
+ hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
+ hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
+ hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
+ }
+}
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/KernelDispatchUtils.h b/iree/compiler/Conversion/LinalgToSPIRV/KernelDispatchUtils.h
index 04be0ff..17f3bc5 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/KernelDispatchUtils.h
+++ b/iree/compiler/Conversion/LinalgToSPIRV/KernelDispatchUtils.h
@@ -32,7 +32,7 @@
namespace mlir {
class FuncOp;
-class LogicalResult;
+struct LogicalResult;
class Operation;
class PatternRewriter;
class ShapedType;
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/Utils.h b/iree/compiler/Conversion/LinalgToSPIRV/Utils.h
index 972e116..ddaecf3 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/Utils.h
+++ b/iree/compiler/Conversion/LinalgToSPIRV/Utils.h
@@ -29,7 +29,7 @@
class SubViewOp;
class OperationFolder;
class OpBuilder;
-class LogicalResult;
+struct LogicalResult;
namespace iree_compiler {
diff --git a/iree/compiler/Dialect/Flow/Transforms/DispatchConfig.cpp b/iree/compiler/Dialect/Flow/Transforms/DispatchConfig.cpp
index 176f09d..9964eae 100644
--- a/iree/compiler/Dialect/Flow/Transforms/DispatchConfig.cpp
+++ b/iree/compiler/Dialect/Flow/Transforms/DispatchConfig.cpp
@@ -31,12 +31,6 @@
namespace {
// TODO(laurenzo): Every one of these should have better support and removed
// from this exclusion list eventually.
-bool isUnsupportedFusionOp(Operation *op) {
- return isa<mhlo::ConcatenateOp, mhlo::ConvOp, mhlo::DotGeneralOp, mhlo::DotOp,
- mhlo::PadOp, mhlo::ReduceOp, mhlo::ReduceWindowOp, mhlo::SliceOp,
- mhlo::TorchIndexSelectOp>(op);
-}
-
// Allowlist of ops that materialize to a an index-permuted copy of some kind
// if they exist standalone. Generally we try to avoid anchoring on these,
// letting them fuse into more meaningful ops as possible.
@@ -182,6 +176,18 @@
return FusionType::DISABLED;
}
+// TODO(b/144530470): replace with tablegen attributes/interfaces.
+bool OpDispatchPolicy::isUnsupportedFusionOp(Operation *op) {
+ return isa<mhlo::ConcatenateOp, mhlo::ConvOp, mhlo::DotGeneralOp, mhlo::DotOp,
+ mhlo::PadOp, mhlo::ReduceOp, mhlo::ReduceWindowOp,
+ mhlo::TorchIndexSelectOp>(op) ||
+ isRootOnlyOp(op);
+}
+
+bool OpDispatchPolicy::isRootOnlyOp(Operation *op) {
+ return isa<mhlo::SliceOp>(op);
+}
+
} // namespace Flow
} // namespace IREE
} // namespace iree_compiler
diff --git a/iree/compiler/Dialect/Flow/Transforms/DispatchConfig.h b/iree/compiler/Dialect/Flow/Transforms/DispatchConfig.h
index 22666b3..ee9299f 100644
--- a/iree/compiler/Dialect/Flow/Transforms/DispatchConfig.h
+++ b/iree/compiler/Dialect/Flow/Transforms/DispatchConfig.h
@@ -41,6 +41,12 @@
OpDispatchPolicy(Dispatchability &dispatchability)
: dispatchability(dispatchability) {}
+ // Returns true if |op| is not able to fuse with either producer or consumer.
+ static bool isUnsupportedFusionOp(Operation *op);
+
+ // Returns true if |op| can only be a root op.
+ static bool isRootOnlyOp(Operation *op);
+
// Returns true if the given |op| can be dispatched in all cases.
// Other passes may handle special cases of these ops but this initial
// identification is conservative.
diff --git a/iree/compiler/Dialect/Flow/Transforms/FoldCompatibleDispatchRegions.cpp b/iree/compiler/Dialect/Flow/Transforms/FoldCompatibleDispatchRegions.cpp
index 984b8c7..9565747 100644
--- a/iree/compiler/Dialect/Flow/Transforms/FoldCompatibleDispatchRegions.cpp
+++ b/iree/compiler/Dialect/Flow/Transforms/FoldCompatibleDispatchRegions.cpp
@@ -13,6 +13,7 @@
// limitations under the License.
#include "iree/compiler/Dialect/Flow/IR/FlowOps.h"
+#include "iree/compiler/Dialect/Flow/Transforms/DispatchConfig.h"
#include "llvm/ADT/ArrayRef.h"
#include "llvm/ADT/DenseMap.h"
#include "llvm/ADT/DenseSet.h"
@@ -200,10 +201,9 @@
// that substituting library calls is easier.
for (auto &block : regionOp.body().getBlocks()) {
for (auto &op : block) {
- // TODO(b/144530470): replace with tablegen attributes/interfaces.
- if (isa<mhlo::ConcatenateOp, mhlo::ConvOp, mhlo::DotGeneralOp,
- mhlo::DotOp, mhlo::PadOp, mhlo::ReduceOp, mhlo::ReduceWindowOp,
- mhlo::SliceOp, mhlo::TorchIndexSelectOp>(op)) {
+ // A root only op is mergable.
+ if (OpDispatchPolicy::isUnsupportedFusionOp(&op) &&
+ !OpDispatchPolicy::isRootOnlyOp(&op)) {
return false;
}
}
@@ -211,6 +211,24 @@
return regionOp.body().getBlocks().size() == 1;
}
+// Returns true if rhs has ops that can only be root op and will lose the
+// characteristic if merge two dispatch regions.
+bool rhsHasRootOnlyOp(DispatchRegionOp &lhs, DispatchRegionOp &rhs) {
+ auto &rhsBlock = rhs.body().front();
+ auto lhsArgs = llvm::to_vector<8>(lhs.args());
+ auto rhsArgs = llvm::to_vector<8>(rhs.args());
+ for (int rhsOpIdx = 0; rhsOpIdx < rhsArgs.size(); ++rhsOpIdx) {
+ for (int lhsResultIdx = 0; lhsResultIdx < lhs.getNumResults();
+ ++lhsResultIdx) {
+ if (rhsArgs[rhsOpIdx] != lhs.getResult(lhsResultIdx)) continue;
+ for (auto *user : rhsBlock.getArgument(rhsOpIdx).getUsers()) {
+ if (OpDispatchPolicy::isRootOnlyOp(user)) return true;
+ }
+ }
+ }
+ return false;
+}
+
// Merges |rhs| into |lhs| and returns the new |lhs| op.
// Precondition: !areDispatchRegionsTransitivelyDependent
DispatchRegionOp mergeDispatchRegions(DispatchRegionOp &lhs,
@@ -345,6 +363,10 @@
LLVM_DEBUG(llvm::dbgs()
<< " -REGION CONTAINS NON-TRIVIAL CONTROL FLOW-\n");
}
+ if (rhsHasRootOnlyOp(lhs, rhs)) {
+ LLVM_DEBUG(llvm::dbgs() << " -RHS REGION HAS ROOT OP-\n");
+ continue;
+ }
mergableRegions[i] = mergeDispatchRegions(lhs, rhs);
if (!mergableRegions[i]) {
return failure();
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/fold_compatible_dispatch_regions.mlir b/iree/compiler/Dialect/Flow/Transforms/test/fold_compatible_dispatch_regions.mlir
index 938c92a..c324ef7 100644
--- a/iree/compiler/Dialect/Flow/Transforms/test/fold_compatible_dispatch_regions.mlir
+++ b/iree/compiler/Dialect/Flow/Transforms/test/fold_compatible_dispatch_regions.mlir
@@ -132,3 +132,33 @@
// CHECK-LABEL: func @dominate
// CHECK: flow.dispatch.region
// CHECK-NOT: flow.dispatch.region
+
+// -----
+
+// Test if the op that only can be a root op fuse with consumer but not
+// producer. This test use a dummy workload to test on root only op
+// functionality.
+module {
+ func @rootOnlyOp(%arg0: tensor<3x4xi32>, %arg1: tensor<1x2xi32>) -> tensor<1x2xi32> {
+ %c0 = constant 0 : index
+ %0 = flow.dispatch.region[%c0 : index](%arg2 = %arg0 : tensor<3x4xi32>) -> tensor<3x4xi32> {
+ %3 = mhlo.add %arg2, %arg2 : tensor<3x4xi32>
+ flow.return %3 : tensor<3x4xi32>
+ }
+ %1 = flow.dispatch.region[%c0 : index](%arg2 = %0 : tensor<3x4xi32>) -> tensor<1x2xi32> {
+ %3 = "mhlo.slice"(%arg2) {limit_indices = dense<[2, 3]> : tensor<2xi64>, start_indices = dense<1> : tensor<2xi64>, strides = dense<1> : tensor<2xi64>} : (tensor<3x4xi32>) -> tensor<1x2xi32>
+ flow.return %3 : tensor<1x2xi32>
+ }
+ %2 = flow.dispatch.region[%c0 : index](%arg2 = %1 : tensor<1x2xi32>, %arg3 = %arg1 : tensor<1x2xi32>) -> tensor<1x2xi32> {
+ %3 = mhlo.multiply %arg2, %arg3 : tensor<1x2xi32>
+ flow.return %3 : tensor<1x2xi32>
+ }
+ return %2 : tensor<1x2xi32>
+ }
+}
+// CHECK-LABEL: func @rootOnlyOp
+// CHECK: flow.dispatch.region
+// CHECK-NEXT: mhlo.add
+// CHECK: flow.dispatch.region
+// CHECK-NEXT: mhlo.slice
+// CHECK-NEXT: mhlo.multiply
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/AOT/BUILD b/iree/compiler/Dialect/HAL/Target/LLVM/AOT/BUILD
index 594a0be..0c81851 100644
--- a/iree/compiler/Dialect/HAL/Target/LLVM/AOT/BUILD
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/AOT/BUILD
@@ -38,7 +38,7 @@
"LLVMAOTTarget.h",
],
deps = [
- ":LLVMAOTTargetLinker",
+ ":LinkerTool",
"//iree/compiler/Dialect/HAL/Target",
"//iree/compiler/Dialect/HAL/Target/LLVM:LLVMBaseTarget",
"//iree/compiler/Dialect/HAL/Target/LLVM:LLVMIRPasses",
@@ -57,16 +57,19 @@
)
cc_library(
- name = "LLVMAOTTargetLinker",
- hdrs = ["LLVMAOTTargetLinker.h"],
- deps = platform_trampoline_deps("LLVMAOTTargetLinker", "compiler/Dialect/HAL/Target/LLVM/AOT"),
+ name = "LinkerTool",
+ srcs = ["LinkerTool.cpp"],
+ hdrs = ["LinkerTool.h"],
+ deps = platform_trampoline_deps("LinkerTools", "compiler/Dialect/HAL/Target/LLVM/AOT"),
)
cc_library(
- name = "LLVMAOTTargetLinker_hdrs",
- hdrs = ["LLVMAOTTargetLinker.h"],
+ name = "LinkerTool_hdrs",
+ hdrs = ["LinkerTool.h"],
deps = [
- "//iree/base:status",
"//iree/compiler/Dialect/HAL/Target/LLVM:LLVMTargetOptions",
+ "@llvm-project//llvm:Core",
+ "@llvm-project//llvm:Support",
+ "@llvm-project//mlir:Support",
],
)
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/AOT/CMakeLists.txt b/iree/compiler/Dialect/HAL/Target/LLVM/AOT/CMakeLists.txt
index ff46a62..7f69d4d 100644
--- a/iree/compiler/Dialect/HAL/Target/LLVM/AOT/CMakeLists.txt
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/AOT/CMakeLists.txt
@@ -26,7 +26,7 @@
SRCS
"LLVMAOTTarget.cpp"
DEPS
- ::LLVMAOTTargetLinker
+ ::LinkerTool
LLVMAArch64AsmParser
LLVMAArch64CodeGen
LLVMARMAsmParser
@@ -46,21 +46,25 @@
iree_cc_library(
NAME
- LLVMAOTTargetLinker
+ LinkerTool
HDRS
- "LLVMAOTTargetLinker.h"
+ "LinkerTool.h"
+ SRCS
+ "LinkerTool.cpp"
DEPS
- iree::compiler::Dialect::HAL::Target::LLVM::AOT::internal::LLVMAOTTargetLinker_internal
+ iree::compiler::Dialect::HAL::Target::LLVM::AOT::internal::LinkerTools_internal
PUBLIC
)
iree_cc_library(
NAME
- LLVMAOTTargetLinker_hdrs
+ LinkerTool_hdrs
HDRS
- "LLVMAOTTargetLinker.h"
+ "LinkerTool.h"
DEPS
- iree::base::status
+ LLVMCore
+ LLVMSupport
+ MLIRSupport
iree::compiler::Dialect::HAL::Target::LLVM::LLVMTargetOptions
PUBLIC
)
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/AOT/LLVMAOTTarget.cpp b/iree/compiler/Dialect/HAL/Target/LLVM/AOT/LLVMAOTTarget.cpp
index 6060600..d86c12c 100644
--- a/iree/compiler/Dialect/HAL/Target/LLVM/AOT/LLVMAOTTarget.cpp
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/AOT/LLVMAOTTarget.cpp
@@ -16,7 +16,7 @@
#include <cstdlib>
-#include "iree/compiler/Dialect/HAL/Target/LLVM/AOT/LLVMAOTTargetLinker.h"
+#include "iree/compiler/Dialect/HAL/Target/LLVM/AOT/LinkerTool.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"
@@ -46,16 +46,20 @@
// 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 (SymbolTable::getSymbolVisibility(func) !=
- SymbolTable::Visibility::Public) {
- nonPublicFn.push_back(func);
- }
- }
- for (auto func : nonPublicFn) {
- func.erase();
+ // We name our files after the executable name so that they are easy to
+ // track both during compilation (logs/artifacts/etc), as outputs (final
+ // intermediate code/binary files), and at runtime (loaded
+ // libraries/symbols/etc).
+ auto libraryName =
+ targetOp.getParentOfType<IREE::HAL::ExecutableOp>().getName().str();
+
+ // TODO(#3737): don't add functions we don't want to serialize to the
+ // module. Right now workgroup count calculation functions end up in here
+ // as std.func ops and not just the llvm.func ops we expect.
+ auto illegalFuncOps =
+ llvm::to_vector<4>(targetOp.getInnerModule().getOps<FuncOp>());
+ for (auto funcOp : illegalFuncOps) {
+ funcOp.erase();
}
// At this moment we are leaving MLIR LLVM dialect land translating module
@@ -63,63 +67,107 @@
auto llvmModule =
mlir::translateModuleToLLVMIR(targetOp.getInnerModule(), context);
if (!llvmModule) {
- return failure();
+ return targetOp.emitError() << "failed to translate the MLIR LLVM "
+ "dialect to the native llvm::Module";
}
+ // Export all entry points such that they are accessible on the dynamic
+ // libraries we generate.
iree::DyLibExecutableDefT dyLibExecutableDef;
- // Create invocation function an populate entry_points.
- auto entryPointOps = targetOp.getBlock().getOps<ExecutableEntryPointOp>();
-
- for (auto entryPointOp : entryPointOps) {
+ SmallVector<StringRef, 8> entryPointNames;
+ for (auto entryPointOp :
+ targetOp.getBlock().getOps<ExecutableEntryPointOp>()) {
dyLibExecutableDef.entry_points.push_back(
std::string(entryPointOp.sym_name()));
+ entryPointNames.push_back(entryPointOp.sym_name());
}
- // LLVMIR opt passes.
+ // Try to grab a linker tool based on the options (and target environment).
+ llvm::Triple targetTriple(options_.targetTriple);
+ auto linkerTool = LinkerTool::getForTarget(targetTriple, options_);
+ if (!linkerTool) {
+ return mlir::emitError(targetOp.getLoc())
+ << "failed to find a target linker for the given target triple '"
+ << options_.targetTriple << "'";
+ }
+
+ // Configure the module with any code generation options required later by
+ // linking (such as initializer functions).
+ if (failed(
+ linkerTool->configureModule(llvmModule.get(), entryPointNames))) {
+ return targetOp.emitError()
+ << "failed to configure LLVM module for target linker";
+ }
+
+ // LLVM opt passes that perform code generation optimizations/transformation
+ // similar to what a frontend would do before passing to linking.
auto targetMachine = createTargetMachine(options_);
if (!targetMachine) {
- targetOp.emitError("Can't create target machine for target triple: " +
- options_.targetTriple);
- return failure();
+ return mlir::emitError(targetOp.getLoc())
+ << "failed to create target machine for target triple '"
+ << options_.targetTriple << "'";
}
-
llvmModule->setDataLayout(targetMachine->createDataLayout());
llvmModule->setTargetTriple(targetMachine->getTargetTriple().str());
-
if (failed(
runLLVMIRPasses(options_, targetMachine.get(), llvmModule.get()))) {
- return targetOp.emitError(
- "Can't build LLVMIR opt passes for ExecutableOp module");
+ return targetOp.emitError()
+ << "failed to run LLVM-IR opt passes for IREE::HAL::ExecutableOp "
+ "targeting '"
+ << options_.targetTriple << "'";
}
- std::string objData;
- if (failed(runEmitObjFilePasses(targetMachine.get(), llvmModule.get(),
- &objData))) {
- return targetOp.emitError("Can't compile LLVMIR module to an obj");
+ // Emit object files.
+ SmallVector<Artifact, 4> objectFiles;
+ {
+ // NOTE: today we just use a single object file, however if we wanted to
+ // scale code generation and linking we'd want to generate one per
+ // function (or something like that).
+ std::string objectData;
+ if (failed(runEmitObjFilePasses(targetMachine.get(), llvmModule.get(),
+ &objectData))) {
+ return targetOp.emitError()
+ << "failed to compile LLVM-IR module to an object file";
+ }
+ auto objectFile = Artifact::createTemporary(libraryName, "obj");
+ auto &os = objectFile.outputFile->os();
+ os << objectData;
+ os.flush();
+ os.close();
+ objectFiles.push_back(std::move(objectFile));
}
- std::string sharedLibData;
- const char *linkerToolPath = std::getenv("IREE_LLVMAOT_LINKER_PATH");
- if (linkerToolPath != nullptr) {
- auto sharedLibDataStatus = linkLLVMAOTObjects(linkerToolPath, objData);
- if (!sharedLibDataStatus.ok()) {
- return targetOp.emitError(
- "Can't link executable and generate target dylib, using linker "
- "toolchain: '" +
- std::string(linkerToolPath) + "'");
- }
- sharedLibData = sharedLibDataStatus.value();
- } else {
- auto sharedLibDataStatus = linkLLVMAOTObjectsWithLLDElf(objData);
- if (!sharedLibDataStatus.ok()) {
- return targetOp.emitError(
- "Can't link executable and generate target dylib using "
- "lld::elf::link");
- }
- sharedLibData = sharedLibDataStatus.value();
+ // Link the generated object files into a dylib.
+ auto linkArtifactsOr =
+ linkerTool->linkDynamicLibrary(libraryName, objectFiles);
+ if (!linkArtifactsOr.hasValue()) {
+ return mlir::emitError(targetOp.getLoc())
+ << "failed to link executable and generate target dylib using "
+ "linker toolchain "
+ << linkerTool->getToolPath();
}
- dyLibExecutableDef.library_embedded = {sharedLibData.begin(),
- sharedLibData.end()};
+ auto &linkArtifacts = linkArtifactsOr.getValue();
+ dyLibExecutableDef.library_embedded =
+ linkArtifacts.libraryFile.read().getValueOr(std::vector<int8_t>());
+ if (dyLibExecutableDef.library_embedded.empty()) {
+ return targetOp.emitError() << "failed to read back dylib temp file at "
+ << linkArtifacts.libraryFile.path;
+ }
+
+ if (options_.debugSymbols && linkArtifacts.debugFile.outputFile) {
+ dyLibExecutableDef.debug_database_embedded =
+ linkArtifacts.debugFile.read().getValue();
+ assert(!dyLibExecutableDef.debug_database_embedded.empty());
+ dyLibExecutableDef.debug_database_filename =
+ llvm::sys::path::filename(linkArtifacts.debugFile.path).str();
+ }
+
+ if (options_.keepLinkerArtifacts) {
+ return mlir::emitRemark(targetOp.getLoc())
+ << "Linker artifacts for " << targetOp.getName() << " preserved:\n"
+ << " " << linkArtifacts.libraryFile.path;
+ linkArtifacts.keepAllFiles();
+ }
::flatbuffers::FlatBufferBuilder fbb;
auto executableOffset =
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/AOT/LLVMAOTTargetLinker.h b/iree/compiler/Dialect/HAL/Target/LLVM/AOT/LLVMAOTTargetLinker.h
deleted file mode 100644
index d6d5220..0000000
--- a/iree/compiler/Dialect/HAL/Target/LLVM/AOT/LLVMAOTTargetLinker.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_COMPILER_DIALECT_HAL_TARGET_LLVM_AOT_LLVMAOTTARGETLINKER_H_
-#define IREE_COMPILER_DIALECT_HAL_TARGET_LLVM_AOT_LLVMAOTTARGETLINKER_H_
-
-#include <string>
-
-#include "iree/base/status.h"
-#include "iree/compiler/Dialect/HAL/Target/LLVM/LLVMTargetOptions.h"
-
-namespace mlir {
-namespace iree_compiler {
-namespace IREE {
-namespace HAL {
-
-// Calls linker tool to link objData and returns shared library blob.
-iree::StatusOr<std::string> linkLLVMAOTObjects(
- const std::string& linkerToolPath, const std::string& objData);
-// Use lld::elf::link for linking objData and returns shared library blob.
-iree::StatusOr<std::string> linkLLVMAOTObjectsWithLLDElf(
- const std::string& objData);
-
-} // namespace HAL
-} // namespace IREE
-} // namespace iree_compiler
-} // namespace mlir
-
-#endif // IREE_COMPILER_DIALECT_HAL_TARGET_LLVM_AOT_LLVMAOTTARGETLINKER_H_
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/AOT/LinkerTool.cpp b/iree/compiler/Dialect/HAL/Target/LLVM/AOT/LinkerTool.cpp
new file mode 100644
index 0000000..083d281
--- /dev/null
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/AOT/LinkerTool.cpp
@@ -0,0 +1,110 @@
+// 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/AOT/LinkerTool.h"
+
+#define DEBUG_TYPE "llvmaot-linker"
+
+namespace mlir {
+namespace iree_compiler {
+namespace IREE {
+namespace HAL {
+
+// static
+Artifact Artifact::createTemporary(StringRef prefix, StringRef suffix) {
+ llvm::SmallString<32> filePath;
+ if (std::error_code error =
+ llvm::sys::fs::createTemporaryFile(prefix, suffix, filePath)) {
+ llvm::errs() << "failed to generate temporary file: " << error.message();
+ return {};
+ }
+ std::error_code error;
+ auto file = std::make_unique<llvm::ToolOutputFile>(filePath, error,
+ llvm::sys::fs::OF_None);
+ if (error) {
+ llvm::errs() << "failed to open temporary file '" << filePath
+ << "': " << error.message();
+ return {};
+ }
+ return {filePath.str().str(), std::move(file)};
+}
+
+// static
+Artifact Artifact::createVariant(StringRef basePath, StringRef suffix) {
+ SmallString<32> filePath(basePath);
+ llvm::sys::path::replace_extension(filePath, suffix);
+ std::error_code error;
+ auto file = std::make_unique<llvm::ToolOutputFile>(filePath, error,
+ llvm::sys::fs::OF_Append);
+ if (error) {
+ llvm::errs() << "failed to open temporary file '" << filePath
+ << "': " << error.message();
+ return {};
+ }
+ return {filePath.str().str(), std::move(file)};
+}
+
+Optional<std::vector<int8_t>> Artifact::read() const {
+ auto fileData = llvm::MemoryBuffer::getFile(path);
+ if (!fileData) {
+ llvm::errs() << "failed to load library output file '" << path << "'";
+ return llvm::None;
+ }
+ auto sourceBuffer = fileData.get()->getBuffer();
+ std::vector<int8_t> resultBuffer(sourceBuffer.size());
+ std::memcpy(resultBuffer.data(), sourceBuffer.data(), sourceBuffer.size());
+ return resultBuffer;
+}
+
+void Artifact::close() { outputFile->os().close(); }
+
+void Artifacts::keepAllFiles() {
+ if (libraryFile.outputFile) libraryFile.outputFile->keep();
+ if (debugFile.outputFile) debugFile.outputFile->keep();
+ for (auto &file : otherFiles) {
+ file.outputFile->keep();
+ }
+}
+
+std::string LinkerTool::getToolPath() const {
+ char *linkerPath = std::getenv("IREE_LLVMAOT_LINKER_PATH");
+ if (linkerPath) {
+ return std::string(linkerPath);
+ } else {
+ return "";
+ }
+}
+
+LogicalResult LinkerTool::runLinkCommand(const std::string &commandLine) {
+ LLVM_DEBUG(llvm::dbgs() << "Running linker command:\n" << commandLine);
+#if defined(_MSC_VER)
+ // It's easy to run afoul of quoting rules on Windows (such as when using
+ // spaces in the linker environment variable). See:
+ // https://stackoverflow.com/a/9965141
+ auto quotedCommandLine = "\"" + commandLine + "\"";
+ int exitCode = system(quotedCommandLine.c_str());
+#else
+ int exitCode = system(commandLine.c_str());
+#endif // _MSC_VER
+ if (exitCode == 0) return success();
+ llvm::errs() << "Linking failed; command line returned exit code " << exitCode
+ << ":\n\n"
+ << commandLine << "\n\n";
+ return failure();
+}
+
+} // namespace HAL
+} // namespace IREE
+} // namespace iree_compiler
+} // namespace mlir
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/AOT/LinkerTool.h b/iree/compiler/Dialect/HAL/Target/LLVM/AOT/LinkerTool.h
new file mode 100644
index 0000000..6f3c39a
--- /dev/null
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/AOT/LinkerTool.h
@@ -0,0 +1,117 @@
+// 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_AOT_LINKERTOOL_H_
+#define IREE_COMPILER_DIALECT_HAL_TARGET_LLVM_AOT_LINKERTOOL_H_
+
+#include <string>
+
+#include "iree/compiler/Dialect/HAL/Target/LLVM/LLVMTargetOptions.h"
+#include "llvm/ADT/Triple.h"
+#include "llvm/IR/Module.h"
+#include "llvm/Support/ToolOutputFile.h"
+#include "mlir/Support/LogicalResult.h"
+
+namespace mlir {
+namespace iree_compiler {
+namespace IREE {
+namespace HAL {
+
+struct Artifact {
+ // Creates an output file path/container pair.
+ // By default the file will be deleted when the link completes; callers must
+ // use llvm::ToolOutputFile::keep() to prevent deletion upon success (or if
+ // leaving artifacts for debugging).
+ static Artifact createTemporary(StringRef prefix, StringRef suffix);
+
+ // Creates an output file derived from the given file's path with a new
+ // suffix.
+ static Artifact createVariant(StringRef basePath, StringRef suffix);
+
+ Artifact() = default;
+ Artifact(std::string path, std::unique_ptr<llvm::ToolOutputFile> outputFile)
+ : path(std::move(path)), outputFile(std::move(outputFile)) {}
+
+ std::string path;
+ std::unique_ptr<llvm::ToolOutputFile> outputFile;
+
+ // Reads the artifact file contents as bytes.
+ Optional<std::vector<int8_t>> read() const;
+
+ // Closes the ostream of the file while preserving the temporary entry on
+ // disk. Use this if files need to be modified by external tools that may
+ // require exclusive access.
+ void close();
+};
+
+struct Artifacts {
+ // File containing the linked library (DLL, ELF, etc).
+ Artifact libraryFile;
+
+ // Optional file containing associated debug information (if stored
+ // separately, such as PDB files).
+ Artifact debugFile;
+
+ // Other files associated with linking.
+ SmallVector<Artifact, 4> otherFiles;
+
+ // Keeps all of the artifacts around after linking completes. Useful for
+ // debugging.
+ void keepAllFiles();
+};
+
+// Base type for linker tools that can turn object files into shared objects.
+class LinkerTool {
+ public:
+ // Gets an instance of a linker tool for the given target options. This may
+ // be a completely different toolchain than that of the host.
+ static std::unique_ptr<LinkerTool> getForTarget(
+ llvm::Triple& targetTriple, LLVMTargetOptions& targetOptions);
+
+ explicit LinkerTool(llvm::Triple targetTriple,
+ LLVMTargetOptions targetOptions)
+ : targetTriple(std::move(targetTriple)),
+ targetOptions(std::move(targetOptions)) {}
+
+ virtual ~LinkerTool() = default;
+
+ // Returns the path to the linker tool binary.
+ virtual std::string getToolPath() const;
+
+ // Configures a module prior to compilation with any additional
+ // functions/exports it may need, such as shared object initializer functions.
+ virtual LogicalResult configureModule(
+ llvm::Module* llvmModule, ArrayRef<StringRef> entryPointNames) = 0;
+
+ // Links the given object files into a dynamically loadable library.
+ // The resulting library (and other associated artifacts) will be returned on
+ // success.
+ virtual Optional<Artifacts> linkDynamicLibrary(
+ StringRef libraryName, ArrayRef<Artifact> objectFiles) = 0;
+
+ protected:
+ // Runs the given command line on the shell, logging failures.
+ LogicalResult runLinkCommand(const std::string& commandLine);
+
+ llvm::Triple targetTriple;
+ LLVMTargetOptions targetOptions;
+};
+
+} // namespace HAL
+} // namespace IREE
+} // namespace iree_compiler
+} // namespace mlir
+
+#endif // IREE_COMPILER_DIALECT_HAL_TARGET_LLVM_AOT_LINKERTOOL_H_
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/AOT/internal/BUILD b/iree/compiler/Dialect/HAL/Target/LLVM/AOT/internal/BUILD
index 6dadc26..bc8a6e0 100644
--- a/iree/compiler/Dialect/HAL/Target/LLVM/AOT/internal/BUILD
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/AOT/internal/BUILD
@@ -19,11 +19,16 @@
)
cc_library(
- name = "LLVMAOTTargetLinker_internal",
- srcs = ["LLVMAOTTargetLinker.cpp"],
+ name = "LinkerTools_internal",
+ srcs = [
+ "LinkerTools.cpp",
+ "UnixLinkerTool.cpp",
+ "WindowsLinkerTool.cpp",
+ ],
deps = [
- "//iree/base:status",
- "//iree/compiler/Dialect/HAL/Target/LLVM/AOT:LLVMAOTTargetLinker_hdrs",
+ "//iree/compiler/Dialect/HAL/Target/LLVM/AOT:LinkerTool_hdrs",
+ "@llvm-project//llvm:Core",
"@llvm-project//llvm:Support",
+ "@llvm-project//mlir:Support",
],
)
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/AOT/internal/CMakeLists.txt b/iree/compiler/Dialect/HAL/Target/LLVM/AOT/internal/CMakeLists.txt
index 01cb459..aafc4a5 100644
--- a/iree/compiler/Dialect/HAL/Target/LLVM/AOT/internal/CMakeLists.txt
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/AOT/internal/CMakeLists.txt
@@ -16,12 +16,15 @@
iree_cc_library(
NAME
- LLVMAOTTargetLinker_internal
+ LinkerTools_internal
SRCS
- "LLVMAOTTargetLinker.cpp"
+ "LinkerTools.cpp"
+ "UnixLinkerTool.cpp"
+ "WindowsLinkerTool.cpp"
DEPS
+ LLVMCore
LLVMSupport
- iree::base::status
- iree::compiler::Dialect::HAL::Target::LLVM::AOT::LLVMAOTTargetLinker_hdrs
+ MLIRSupport
+ iree::compiler::Dialect::HAL::Target::LLVM::AOT::LinkerTool_hdrs
PUBLIC
)
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/AOT/internal/LLVMAOTTargetLinker.cpp b/iree/compiler/Dialect/HAL/Target/LLVM/AOT/internal/LLVMAOTTargetLinker.cpp
deleted file mode 100644
index 8acd2d6..0000000
--- a/iree/compiler/Dialect/HAL/Target/LLVM/AOT/internal/LLVMAOTTargetLinker.cpp
+++ /dev/null
@@ -1,79 +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/AOT/LLVMAOTTargetLinker.h"
-
-#include "iree/base/status.h"
-#include "llvm/Support/ToolOutputFile.h"
-
-namespace mlir {
-namespace iree_compiler {
-namespace IREE {
-namespace HAL {
-
-iree::StatusOr<std::string> linkLLVMAOTObjects(
- const std::string& linkerToolPath, const std::string& objData) {
- llvm::SmallString<32> objFilePath, dylibFilePath;
- if (std::error_code error = llvm::sys::fs::createTemporaryFile(
- "llvmaot_dylibs", "objfile", objFilePath)) {
- return iree::InternalErrorBuilder(IREE_LOC)
- << "Failed to generate temporary file for objfile : '"
- << error.message() << "'";
- }
- if (std::error_code error = llvm::sys::fs::createTemporaryFile(
- "llvmaot_dylibs", "dylibfile", dylibFilePath)) {
- return iree::InternalErrorBuilder(IREE_LOC)
- << "Failed to generate temporary file for dylib : '"
- << error.message() << "'";
- }
- std::error_code error;
- auto outputFile = std::make_unique<llvm::ToolOutputFile>(
- objFilePath, error, llvm::sys::fs::F_None);
- if (error) {
- return iree::InternalErrorBuilder(IREE_LOC)
- << "Failed to open temporary objfile '" << objFilePath.c_str()
- << "' for dylib : '" << error.message() << "'";
- }
-
- outputFile->os() << objData;
- outputFile->os().flush();
-
- auto linkingCmd =
- (linkerToolPath + " -shared " + objFilePath + " -o " + dylibFilePath)
- .str();
- int systemRet = system(linkingCmd.c_str());
- if (systemRet != 0) {
- return iree::InternalErrorBuilder(IREE_LOC)
- << linkingCmd << " failed with exit code " << systemRet;
- }
-
- auto dylibData = llvm::MemoryBuffer::getFile(dylibFilePath);
- if (!dylibData) {
- return iree::InternalErrorBuilder(IREE_LOC)
- << "Failed to read temporary dylib file '" << dylibFilePath.c_str()
- << "'";
- }
- return dylibData.get()->getBuffer().str();
-}
-
-iree::StatusOr<std::string> linkLLVMAOTObjectsWithLLDElf(
- const std::string& objData) {
- return iree::UnimplementedErrorBuilder(IREE_LOC)
- << "linkLLVMAOTObjectsWithLLD not implemented yet!";
-}
-
-} // namespace HAL
-} // namespace IREE
-} // namespace iree_compiler
-} // namespace mlir
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/AOT/internal/LinkerTools.cpp b/iree/compiler/Dialect/HAL/Target/LLVM/AOT/internal/LinkerTools.cpp
new file mode 100644
index 0000000..6fc4b93
--- /dev/null
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/AOT/internal/LinkerTools.cpp
@@ -0,0 +1,43 @@
+// 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/AOT/LinkerTool.h"
+
+namespace mlir {
+namespace iree_compiler {
+namespace IREE {
+namespace HAL {
+
+// TODO(benvanik): add other platforms:
+// createMacLinkerTool using ld64.lld
+// createWasmLinkerTool wasm-ld
+
+std::unique_ptr<LinkerTool> createUnixLinkerTool(
+ llvm::Triple &targetTriple, LLVMTargetOptions &targetOptions);
+std::unique_ptr<LinkerTool> createWindowsLinkerTool(
+ llvm::Triple &targetTriple, LLVMTargetOptions &targetOptions);
+
+// static
+std::unique_ptr<LinkerTool> LinkerTool::getForTarget(
+ llvm::Triple &targetTriple, LLVMTargetOptions &targetOptions) {
+ if (targetTriple.isOSWindows() || targetTriple.isWindowsMSVCEnvironment()) {
+ return createWindowsLinkerTool(targetTriple, targetOptions);
+ }
+ return createUnixLinkerTool(targetTriple, targetOptions);
+}
+
+} // namespace HAL
+} // namespace IREE
+} // namespace iree_compiler
+} // namespace mlir
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/AOT/internal/UnixLinkerTool.cpp b/iree/compiler/Dialect/HAL/Target/LLVM/AOT/internal/UnixLinkerTool.cpp
new file mode 100644
index 0000000..fcd1986
--- /dev/null
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/AOT/internal/UnixLinkerTool.cpp
@@ -0,0 +1,86 @@
+// 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/AOT/LinkerTool.h"
+#include "llvm/IR/Function.h"
+#include "llvm/IR/IRBuilder.h"
+#include "llvm/Support/FormatVariadic.h"
+
+#define DEBUG_TYPE "llvmaot-linker"
+
+namespace mlir {
+namespace iree_compiler {
+namespace IREE {
+namespace HAL {
+
+// Unix linker (ld-like); for ELF files.
+class UnixLinkerTool : public LinkerTool {
+ public:
+ using LinkerTool::LinkerTool;
+
+ std::string getToolPath() const override {
+ auto toolPath = LinkerTool::getToolPath();
+ return toolPath.empty() ? "ld.lld" : toolPath;
+ }
+
+ LogicalResult configureModule(llvm::Module *llvmModule,
+ ArrayRef<StringRef> entryPointNames) override {
+ // Possibly a no-op in ELF files; needs to be verified.
+ return success();
+ }
+
+ Optional<Artifacts> linkDynamicLibrary(
+ StringRef libraryName, ArrayRef<Artifact> objectFiles) override {
+ Artifacts artifacts;
+
+ // Create the shared object name; if we only have a single input object we
+ // can just reuse that.
+ if (objectFiles.size() == 1) {
+ artifacts.libraryFile =
+ Artifact::createVariant(objectFiles.front().path, "so");
+ } else {
+ artifacts.libraryFile = Artifact::createTemporary(libraryName, "so");
+ }
+ artifacts.libraryFile.close();
+
+ SmallVector<std::string, 8> flags = {
+ getToolPath(),
+ "-shared",
+ "-o " + artifacts.libraryFile.path,
+ };
+
+ // TODO(ataei): add flags based on targetTriple.isAndroid(), like
+ // -static-libstdc++ (if this is needed, which it shouldn't be).
+
+ // Link all input objects. Note that we are not linking whole-archive as we
+ // want to allow dropping of unused codegen outputs.
+ for (auto &objectFile : objectFiles) {
+ flags.push_back(objectFile.path);
+ }
+
+ auto commandLine = llvm::join(flags, " ");
+ if (failed(runLinkCommand(commandLine))) return llvm::None;
+ return artifacts;
+ }
+};
+
+std::unique_ptr<LinkerTool> createUnixLinkerTool(
+ llvm::Triple &targetTriple, LLVMTargetOptions &targetOptions) {
+ return std::make_unique<UnixLinkerTool>(targetTriple, targetOptions);
+}
+
+} // namespace HAL
+} // namespace IREE
+} // namespace iree_compiler
+} // namespace mlir
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/AOT/internal/WindowsLinkerTool.cpp b/iree/compiler/Dialect/HAL/Target/LLVM/AOT/internal/WindowsLinkerTool.cpp
new file mode 100644
index 0000000..40632e3
--- /dev/null
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/AOT/internal/WindowsLinkerTool.cpp
@@ -0,0 +1,290 @@
+// 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/AOT/LinkerTool.h"
+#include "llvm/IR/Function.h"
+#include "llvm/IR/IRBuilder.h"
+#include "llvm/Support/FormatVariadic.h"
+
+#define DEBUG_TYPE "llvmaot-linker"
+
+namespace mlir {
+namespace iree_compiler {
+namespace IREE {
+namespace HAL {
+
+// Windows linker (MSVC link.exe-like); for DLL files.
+class WindowsLinkerTool : public LinkerTool {
+ public:
+ using LinkerTool::LinkerTool;
+
+ std::string getToolPath() const override {
+ auto toolPath = LinkerTool::getToolPath();
+ return toolPath.empty() ? "lld-link" : toolPath;
+ }
+
+ LogicalResult configureModule(llvm::Module *llvmModule,
+ ArrayRef<StringRef> entryPointNames) override {
+ auto &ctx = llvmModule->getContext();
+
+ // Create a _DllMainCRTStartup replacement that does not initialize the CRT.
+ // This is required to prevent a bunch of CRT junk (locale, errno, TLS, etc)
+ // from getting emitted in such a way that it cannot be stripped by LTCG.
+ // Since we don't emit code using the CRT (beyond memset/memcpy) this is
+ // fine and can reduce binary sizes by 50-100KB.
+ //
+ // More info:
+ // https://docs.microsoft.com/en-us/cpp/build/run-time-library-behavior?view=vs-2019
+ {
+ auto dwordType = llvm::IntegerType::get(ctx, 32);
+ auto ptrType = llvm::PointerType::getUnqual(dwordType);
+ auto entry = cast<llvm::Function>(
+ llvmModule
+ ->getOrInsertFunction("iree_dll_main", dwordType, ptrType,
+ dwordType, ptrType)
+ .getCallee());
+ entry->setCallingConv(llvm::CallingConv::X86_StdCall);
+ entry->setDLLStorageClass(
+ llvm::GlobalValue::DLLStorageClassTypes::DLLExportStorageClass);
+ entry->setLinkage(llvm::GlobalValue::LinkageTypes::ExternalLinkage);
+ auto *block = llvm::BasicBlock::Create(ctx, "entry", entry);
+ llvm::IRBuilder<> builder(block);
+ auto one = llvm::ConstantInt::get(dwordType, 1, false);
+ builder.CreateRet(one);
+ }
+
+ // For now we ensure that our entry points are exported (via linker
+ // directives embedded in the object file) and in a compatible calling
+ // convention.
+ // TODO(benvanik): switch to executable libraries w/ internal functions.
+ for (auto entryPointName : entryPointNames) {
+ auto *entryPointFn = llvmModule->getFunction(entryPointName);
+ entryPointFn->setCallingConv(llvm::CallingConv::X86_StdCall);
+ entryPointFn->setDLLStorageClass(
+ llvm::GlobalValue::DLLStorageClassTypes::DLLExportStorageClass);
+ entryPointFn->setLinkage(
+ llvm::GlobalValue::LinkageTypes::ExternalLinkage);
+ entryPointFn->setVisibility(
+ llvm::GlobalValue::VisibilityTypes::DefaultVisibility);
+ entryPointFn->addFnAttr(llvm::Attribute::UWTable);
+ }
+
+ return success();
+ }
+
+ Optional<Artifacts> linkDynamicLibrary(
+ StringRef libraryName, ArrayRef<Artifact> objectFiles) override {
+ Artifacts artifacts;
+
+ // Create the shared object name; if we only have a single input object we
+ // can just reuse that.
+ if (objectFiles.size() == 1) {
+ artifacts.libraryFile =
+ Artifact::createVariant(objectFiles.front().path, "dll");
+ } else {
+ artifacts.libraryFile = Artifact::createTemporary(libraryName, "dll");
+ }
+
+ // link.exe doesn't like the files being opened. We don't use them as
+ // streams so close them all now before running the linker.
+ artifacts.libraryFile.close();
+
+ // We need a full path for the PDB and I hate strings in LLVM grumble.
+ SmallString<32> pdbPath(artifacts.libraryFile.path);
+ llvm::sys::path::replace_extension(pdbPath, "pdb");
+
+ SmallVector<std::string, 8> flags = {
+ getToolPath(),
+
+ // Useful when debugging linking/loading issues:
+ // "/verbose",
+
+ // https://docs.microsoft.com/en-us/cpp/build/reference/dll-build-a-dll?view=vs-2019
+ // Builds a DLL and exports functions with the dllexport storage class.
+ "/dll",
+
+ // Forces a fixed timestamp to ensure files are reproducable across
+ // builds. Undocumented but accepted by both link and lld-link.
+ // https://blog.conan.io/2019/09/02/Deterministic-builds-with-C-C++.html
+ "/Brepro",
+
+ // https://docs.microsoft.com/en-us/cpp/build/reference/nodefaultlib-ignore-libraries?view=vs-2019
+ // Ignore any libraries that are specified by the platform as we
+ // directly provide the ones we want.
+ "/nodefaultlib",
+
+ // https://docs.microsoft.com/en-us/cpp/build/reference/incremental-link-incrementally?view=vs-2019
+ // Disable incremental linking as we are only ever linking in one-shot
+ // mode to temp files. This avoids additional file padding and ordering
+ // restrictions that enable incremental linking. Our other options will
+ // prevent incremental linking in most cases, but it doesn't hurt to be
+ // explicit.
+ "/incremental:no",
+
+ // https://docs.microsoft.com/en-us/cpp/build/reference/guard-enable-guard-checks?view=vs-2019
+ // No control flow guard lookup (indirect branch verification).
+ "/guard:no",
+
+ // https://docs.microsoft.com/en-us/cpp/build/reference/safeseh-image-has-safe-exception-handlers?view=vs-2019
+ // We don't want exception unwind tables in our output.
+ "/safeseh:no",
+
+ // https://docs.microsoft.com/en-us/cpp/build/reference/entry-entry-point-symbol?view=vs-2019
+ // Use our entry point instead of the standard CRT one; ensures that we
+ // pull in no global state from the CRT.
+ "/entry:iree_dll_main",
+
+ // https://docs.microsoft.com/en-us/cpp/build/reference/debug-generate-debug-info?view=vs-2019
+ // Copies all PDB information into the final PDB so that we can use the
+ // same PDB across multiple machines.
+ "/debug:full",
+
+ // https://docs.microsoft.com/en-us/cpp/build/reference/pdb-use-program-database
+ // Generates the PDB file containing the debug information.
+ ("/pdb:" + pdbPath).str(),
+
+ // https://docs.microsoft.com/en-us/cpp/build/reference/pdbaltpath-use-alternate-pdb-path?view=vs-2019
+ // Forces the PDB we generate to be referenced in the DLL as just a
+ // relative path to the DLL itself. This allows us to move the PDBs
+ // along with the build DLLs across machines.
+ "/pdbaltpath:%_PDB%",
+
+ // https://docs.microsoft.com/en-us/cpp/build/reference/out-output-file-name?view=vs-2019
+ // Target for linker output. The base name of this path will be used for
+ // additional output files (like the map and pdb).
+ "/out:" + artifacts.libraryFile.path,
+ };
+
+ if (targetOptions.optLevel.getSpeedupLevel() >= 2 ||
+ targetOptions.optLevel.getSizeLevel() >= 2) {
+ // https://docs.microsoft.com/en-us/cpp/build/reference/opt-optimizations?view=vs-2019
+ // Enable all the fancy optimizations.
+ flags.push_back("/opt:ref,icf,lbr");
+ }
+
+ // SDK and MSVC paths.
+ // These rely on the environment variables provided by the
+ // vcvarsall or VsDevCmd ("Developer Command Prompt") scripts. They can also
+ // be manually be specified.
+ //
+ // We could also check to see if vswhere is installed and query that in the
+ // event of missing environment variables; that would eliminate the need for
+ // specifying things from for example IDEs that may not bring in the vcvars.
+ //
+ /* Example values:
+ UCRTVersion=10.0.18362.0
+ UniversalCRTSdkDir=C:\Program Files (x86)\Windows Kits\10\
+ VCToolsInstallDir=C:\Program Files (x86)\Microsoft Visual
+ Studio\2019\Preview\VC\Tools\MSVC\14.28.29304\
+ */
+ if (!getenv("VCToolsInstallDir") || !getenv("UniversalCRTSdkDir")) {
+ llvm::errs() << "required environment for lld-link/link not specified; "
+ "ensure you are building from a shell where "
+ "vcvarsall/VsDevCmd.bat/etc has been used";
+ return llvm::None;
+ }
+ const char *arch;
+ if (targetTriple.isARM() && targetTriple.isArch32Bit()) {
+ arch = "arm";
+ } else if (targetTriple.isARM()) {
+ arch = "arm64";
+ } else if (targetTriple.isX86() && targetTriple.isArch32Bit()) {
+ arch = "x86";
+ } else if (targetTriple.isX86()) {
+ arch = "x64";
+ } else {
+ llvm::errs() << "unsupported Windows target triple (no arch libs): "
+ << targetTriple.str();
+ return llvm::None;
+ }
+ flags.push_back(
+ llvm::formatv("/libpath:\"{0}\\lib\\{1}\"", "%VCToolsInstallDir%", arch)
+ .str());
+ flags.push_back(llvm::formatv("/libpath:\"{0}\\Lib\\{1}\\ucrt\\{2}\"",
+ "%UniversalCRTSdkDir%", "%UCRTVersion%", arch)
+ .str());
+ flags.push_back(llvm::formatv("/libpath:\"{0}\\Lib\\{1}\\um\\{2}\"",
+ "%UniversalCRTSdkDir%", "%UCRTVersion%", arch)
+ .str());
+
+ // We need to link against different libraries based on our configuration
+ // matrix (dynamic/static and debug/release).
+ int libIndex = 0;
+ if (targetOptions.optLevel.getSpeedupLevel() == 0) {
+ libIndex += 0; // debug
+ } else {
+ libIndex += 2; // release
+ }
+ libIndex += targetOptions.linkStatic ? 1 : 0;
+
+ // The required libraries for linking DLLs:
+ // https://docs.microsoft.com/en-us/cpp/c-runtime-library/crt-library-features?view=msvc-160
+ //
+ // NOTE: there are only static versions of msvcrt as it's the startup code.
+ static const char *kMSVCRTLibs[4] = {
+ /* debug/dynamic */ "msvcrtd.lib",
+ /* debug/static */ "msvcrtd.lib",
+ /* release/dynamic */ "msvcrt.lib",
+ /* release/static */ "msvcrt.lib",
+ };
+ static const char *kVCRuntimeLibs[4] = {
+ /* debug/dynamic */ "vcruntimed.lib",
+ /* debug/static */ "libvcruntimed.lib",
+ /* release/dynamic */ "vcruntime.lib",
+ /* release/static */ "libvcruntime.lib",
+ };
+ static const char *kUCRTLibs[4] = {
+ /* debug/dynamic */ "ucrtd.lib",
+ /* debug/static */ "libucrtd.lib",
+ /* release/dynamic */ "ucrt.lib",
+ /* release/static */ "libucrt.lib",
+ };
+ flags.push_back(kMSVCRTLibs[libIndex]);
+ flags.push_back(kVCRuntimeLibs[libIndex]);
+ flags.push_back(kUCRTLibs[libIndex]);
+ flags.push_back("kernel32.lib");
+
+ // Link all input objects. Note that we are not linking whole-archive as we
+ // want to allow dropping of unused codegen outputs.
+ for (auto &objectFile : objectFiles) {
+ flags.push_back(objectFile.path);
+ }
+
+ auto commandLine = llvm::join(flags, " ");
+ if (failed(runLinkCommand(commandLine))) return llvm::None;
+
+ // PDB file gets generated wtih the same path + .pdb.
+ artifacts.debugFile =
+ Artifact::createVariant(artifacts.libraryFile.path, "pdb");
+
+ // We currently discard some of the other file outputs (like the .exp
+ // listing the exported symbols) as we don't need them.
+ artifacts.otherFiles.push_back(
+ Artifact::createVariant(artifacts.libraryFile.path, "exp"));
+ artifacts.otherFiles.push_back(
+ Artifact::createVariant(artifacts.libraryFile.path, "lib"));
+
+ return artifacts;
+ }
+};
+
+std::unique_ptr<LinkerTool> createWindowsLinkerTool(
+ llvm::Triple &targetTriple, LLVMTargetOptions &targetOptions) {
+ return std::make_unique<WindowsLinkerTool>(targetTriple, targetOptions);
+}
+
+} // namespace HAL
+} // namespace IREE
+} // namespace iree_compiler
+} // namespace mlir
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/BUILD b/iree/compiler/Dialect/HAL/Target/LLVM/BUILD
index 9c608c2..9040b96 100644
--- a/iree/compiler/Dialect/HAL/Target/LLVM/BUILD
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/BUILD
@@ -80,6 +80,7 @@
"LLVMTargetOptions.h",
],
deps = [
+ "@llvm-project//llvm:MC",
"@llvm-project//llvm:Passes",
"@llvm-project//llvm:Support",
"@llvm-project//llvm:Target",
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/CMakeLists.txt b/iree/compiler/Dialect/HAL/Target/LLVM/CMakeLists.txt
index 1f0f77f..416f32e 100644
--- a/iree/compiler/Dialect/HAL/Target/LLVM/CMakeLists.txt
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/CMakeLists.txt
@@ -66,6 +66,7 @@
SRCS
"LLVMTargetOptions.cpp"
DEPS
+ LLVMMC
LLVMPasses
LLVMSupport
LLVMTarget
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/LLVMBaseTarget.cpp b/iree/compiler/Dialect/HAL/Target/LLVM/LLVMBaseTarget.cpp
index 8efcf6c..4861dd3 100644
--- a/iree/compiler/Dialect/HAL/Target/LLVM/LLVMBaseTarget.cpp
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/LLVMBaseTarget.cpp
@@ -80,13 +80,37 @@
buildLLVMTransformPassPipeline(passManager);
}
+static FileLineColLoc findFirstFileLoc(Location baseLoc) {
+ if (auto loc = baseLoc.dyn_cast<FusedLoc>()) {
+ for (auto &childLoc : loc.getLocations()) {
+ auto childResult = findFirstFileLoc(childLoc);
+ if (childResult) return childResult;
+ }
+ } else if (auto loc = baseLoc.dyn_cast<FileLineColLoc>()) {
+ return loc;
+ }
+ return FileLineColLoc{};
+}
+
+static std::string guessModuleName(mlir::ModuleOp moduleOp) {
+ std::string moduleName =
+ moduleOp.getName().hasValue() ? moduleOp.getName().getValue().str() : "";
+ if (!moduleName.empty()) return moduleName;
+ FileLineColLoc loc = findFirstFileLoc(moduleOp.getLoc());
+ return llvm::sys::path::stem(loc.getFilename()).str();
+}
+
LogicalResult LLVMBaseTargetBackend::linkExecutables(mlir::ModuleOp moduleOp) {
OpBuilder builder = OpBuilder::atBlockBegin(moduleOp.getBody());
auto executableOps =
llvm::to_vector<8>(moduleOp.getOps<IREE::HAL::ExecutableOp>());
+ // Guess a module name, if needed, to make the output files readable.
+ auto moduleName = guessModuleName(moduleOp);
+
// Create our new "linked" hal.executable.
- std::string linkedExecutableName = llvm::formatv("linked_{0}", name());
+ std::string linkedExecutableName =
+ llvm::formatv("{0}_linked_{1}", moduleName, name());
auto linkedExecutableOp = builder.create<IREE::HAL::ExecutableOp>(
moduleOp.getLoc(), linkedExecutableName);
linkedExecutableOp.setPrivate();
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/LLVMIRPasses.cpp b/iree/compiler/Dialect/HAL/Target/LLVM/LLVMIRPasses.cpp
index de66446..65cd442 100644
--- a/iree/compiler/Dialect/HAL/Target/LLVM/LLVMIRPasses.cpp
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/LLVMIRPasses.cpp
@@ -31,16 +31,33 @@
namespace IREE {
namespace HAL {
+static llvm::CodeGenOpt::Level passBuilderOptLevelToCodeGenOptLevel(
+ const llvm::PassBuilder::OptimizationLevel &level) {
+ switch (level.getSpeedupLevel()) {
+ case 0:
+ return llvm::CodeGenOpt::None;
+ case 1:
+ return llvm::CodeGenOpt::Less;
+ case 2:
+ default:
+ return llvm::CodeGenOpt::Default;
+ case 3:
+ return llvm::CodeGenOpt::Aggressive;
+ }
+}
+
std::unique_ptr<llvm::TargetMachine> createTargetMachine(
const LLVMTargetOptions &targetOptions) {
std::string errorMessage;
auto target = llvm::TargetRegistry::lookupTarget(targetOptions.targetTriple,
errorMessage);
if (!target) return nullptr;
- // TODO(ataei): Once we have an AOT backend pass cpu and cpu-features
std::unique_ptr<llvm::TargetMachine> machine(target->createTargetMachine(
- targetOptions.targetTriple, "generic" /* cpu e.g k8*/,
- "" /* cpu features e.g avx512fma*/, targetOptions.options, {}));
+ targetOptions.targetTriple, targetOptions.targetCPU /* cpu e.g k8*/,
+ targetOptions.targetCPUFeatures /* cpu features e.g avx512fma*/,
+ targetOptions.options, {}, {},
+ passBuilderOptLevelToCodeGenOptLevel(targetOptions.optLevel),
+ /*JIT=*/false));
return machine;
}
@@ -68,10 +85,12 @@
passBuilder.registerLoopAnalyses(loopAnalysisManager);
passBuilder.crossRegisterProxies(loopAnalysisManager, functionAnalysisManager,
cGSCCAnalysisManager, moduleAnalysisManager);
- llvm::ModulePassManager modulePassManager;
- modulePassManager =
- passBuilder.buildPerModuleDefaultPipeline(options.optLevel);
- modulePassManager.run(*module, moduleAnalysisManager);
+ if (options.optLevel != llvm::PassBuilder::OptimizationLevel::O0) {
+ llvm::ModulePassManager modulePassManager;
+ modulePassManager =
+ passBuilder.buildPerModuleDefaultPipeline(options.optLevel);
+ modulePassManager.run(*module, moduleAnalysisManager);
+ }
if (llvm::verifyModule(*module)) return failure();
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/LLVMTargetOptions.cpp b/iree/compiler/Dialect/HAL/Target/LLVM/LLVMTargetOptions.cpp
index e71beac..649dbf0 100644
--- a/iree/compiler/Dialect/HAL/Target/LLVM/LLVMTargetOptions.cpp
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/LLVMTargetOptions.cpp
@@ -15,6 +15,7 @@
#include "iree/compiler/Dialect/HAL/Target/LLVM/LLVMTargetOptions.h"
#include "llvm/ADT/APFloat.h"
+#include "llvm/MC/SubtargetFeature.h"
#include "llvm/Support/CommandLine.h"
#include "llvm/Support/Host.h"
#include "llvm/Target/TargetOptions.h"
@@ -26,17 +27,34 @@
LLVMTargetOptions getDefaultLLVMTargetOptions() {
LLVMTargetOptions targetOptions;
+
// Host target triple.
targetOptions.targetTriple = llvm::sys::getDefaultTargetTriple();
+ targetOptions.targetCPU = llvm::sys::getHostCPUName().str();
+ {
+ llvm::SubtargetFeatures features;
+ llvm::StringMap<bool> hostFeatures;
+ if (llvm::sys::getHostCPUFeatures(hostFeatures)) {
+ for (auto &feature : hostFeatures) {
+ features.AddFeature(feature.first(), feature.second);
+ }
+ }
+ targetOptions.targetCPUFeatures = features.getString();
+ }
+
// LLVM loop optimization options.
targetOptions.pipelineTuningOptions.LoopInterleaving = true;
targetOptions.pipelineTuningOptions.LoopVectorization = true;
targetOptions.pipelineTuningOptions.LoopUnrolling = true;
+
// LLVM SLP Auto vectorizer.
targetOptions.pipelineTuningOptions.SLPVectorization = true;
+
// LLVM -O3.
+ // TODO(benvanik): add an option for this.
targetOptions.optLevel = llvm::PassBuilder::OptimizationLevel::O3;
targetOptions.options.FloatABIType = llvm::FloatABI::Hard;
+
return targetOptions;
}
@@ -46,16 +64,57 @@
static llvm::cl::opt<std::string> clTargetTriple(
"iree-llvm-target-triple", llvm::cl::desc("LLVM target machine triple"),
llvm::cl::init(llvmTargetOptions.targetTriple));
- static llvm::cl::opt<bool> clSoftFloat(
- "iree-llvm-enable-msoft-float-abi",
+ static llvm::cl::opt<std::string> clTargetCPU(
+ "iree-llvm-target-cpu",
+ llvm::cl::desc(
+ "LLVM target machine CPU; use 'host' for your host native CPU"),
+ llvm::cl::init("generic"));
+ static llvm::cl::opt<std::string> clTargetCPUFeatures(
+ "iree-llvm-target-cpu-features",
+ llvm::cl::desc("LLVM target machine CPU features; use 'host' for your "
+ "host native CPU"),
+ llvm::cl::init(""));
+ llvmTargetOptions.targetTriple = clTargetTriple;
+ if (clTargetCPU != "host") {
+ llvmTargetOptions.targetCPU = clTargetCPU;
+ }
+ if (clTargetCPUFeatures != "host") {
+ llvmTargetOptions.targetCPUFeatures = clTargetCPUFeatures;
+ }
+
+ static llvm::cl::opt<llvm::FloatABI::ABIType> clTargetFloatABI(
+ "iree-llvm-target-float-abi",
llvm::cl::desc("LLVM target codegen enables soft float abi e.g "
"-mfloat-abi=softfp"),
- llvm::cl::init(false));
+ llvm::cl::init(llvmTargetOptions.options.FloatABIType),
+ llvm::cl::values(
+ clEnumValN(llvm::FloatABI::Default, "default", "Default (softfp)"),
+ clEnumValN(llvm::FloatABI::Soft, "soft",
+ "Software floating-point emulation"),
+ clEnumValN(llvm::FloatABI::Hard, "hard",
+ "Hardware floating-point instructions")));
+ llvmTargetOptions.options.FloatABIType = clTargetFloatABI;
- llvmTargetOptions.targetTriple = clTargetTriple;
- if (clSoftFloat) {
- llvmTargetOptions.options.FloatABIType = llvm::FloatABI::Soft;
- }
+ static llvm::cl::opt<bool> clDebugSymbols(
+ "iree-llvm-debug-symbols",
+ llvm::cl::desc("Generate and embed debug information (DWARF, PDB, etc)"),
+ llvm::cl::init(llvmTargetOptions.debugSymbols));
+ llvmTargetOptions.debugSymbols = clDebugSymbols;
+
+ static llvm::cl::opt<bool> clLinkStatic(
+ "iree-llvm-link-static",
+ llvm::cl::desc(
+ "Links system libraries into binaries statically to isolate them "
+ "from platform dependencies needed at runtime"),
+ llvm::cl::init(llvmTargetOptions.linkStatic));
+ llvmTargetOptions.linkStatic = clLinkStatic;
+
+ static llvm::cl::opt<bool> clKeepLinkerArtifacts(
+ "iree-llvm-keep-linker-artifacts",
+ llvm::cl::desc("Keep LLVM linker target artifacts (.so/.dll/etc)"),
+ llvm::cl::init(llvmTargetOptions.keepLinkerArtifacts));
+ llvmTargetOptions.keepLinkerArtifacts = clKeepLinkerArtifacts;
+
return llvmTargetOptions;
}
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/LLVMTargetOptions.h b/iree/compiler/Dialect/HAL/Target/LLVM/LLVMTargetOptions.h
index 4893566..ba454b6 100644
--- a/iree/compiler/Dialect/HAL/Target/LLVM/LLVMTargetOptions.h
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/LLVMTargetOptions.h
@@ -24,10 +24,28 @@
namespace HAL {
struct LLVMTargetOptions {
+ // Target machine configuration.
+ std::string targetTriple;
+ std::string targetCPU;
+ std::string targetCPUFeatures;
+
llvm::PipelineTuningOptions pipelineTuningOptions;
llvm::PassBuilder::OptimizationLevel optLevel;
llvm::TargetOptions options;
- std::string targetTriple;
+
+ // Include debug information in output files (PDB, DWARF, etc).
+ // Though this can be set independently from the optLevel (so -O3 with debug
+ // information is valid) it may significantly change the output program
+ // and benchmarking
+ bool debugSymbols = true;
+
+ // Link any required runtime libraries into the produced binaries statically.
+ // This increases resulting binary size but enables the binaries to be used on
+ // any machine without requiring matching system libraries to be installed.
+ bool linkStatic = false;
+
+ // True to keep linker artifacts for debugging.
+ bool keepLinkerArtifacts = false;
};
// Returns LLVMTargetOptions struct intialized with the
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/test/binaryop_test.mlir b/iree/compiler/Dialect/HAL/Target/LLVM/test/binary_op.mlir
similarity index 90%
rename from iree/compiler/Dialect/HAL/Target/LLVM/test/binaryop_test.mlir
rename to iree/compiler/Dialect/HAL/Target/LLVM/test/binary_op.mlir
index 0a66d94..41f1939 100644
--- a/iree/compiler/Dialect/HAL/Target/LLVM/test/binaryop_test.mlir
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/test/binary_op.mlir
@@ -11,7 +11,7 @@
}
}
-// CHECK-LABEL: hal.executable @linked_llvm_ir
+// CHECK-LABEL: hal.executable @binary_op_linked_llvm_ir
// CHECK-DAG: hal.executable.binary attributes {
// CHECK-SAME: data = dense
// CHECK-SAME: format = 1280071245 : 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 6744a70..1a48518 100644
--- a/iree/compiler/Dialect/HAL/Target/LLVM/test/matmul_op.mlir
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/test/matmul_op.mlir
@@ -11,7 +11,7 @@
}
}
-// CHECK-LABEL: hal.executable @linked_llvm_ir
+// CHECK-LABEL: hal.executable @matmul_op_linked_llvm_ir
// CHECK-DAG: hal.executable.binary attributes {
// CHECK-SAME: data = dense
// CHECK-SAME: format = 1280071245 : i32} {
diff --git a/iree/hal/dylib/dylib_executable.cc b/iree/hal/dylib/dylib_executable.cc
index f3d79cf..3d3ff01 100644
--- a/iree/hal/dylib/dylib_executable.cc
+++ b/iree/hal/dylib/dylib_executable.cc
@@ -34,9 +34,13 @@
DyLibExecutable::~DyLibExecutable() {
IREE_TRACE_SCOPE0("DyLibExecutable::dtor");
- // TODO(benvanik): move to an atexit handler when tracing is enabled.
- // executable_library_.release();
+#if IREE_TRACING_FEATURES & IREE_TRACING_FEATURE_INSTRUMENTATION
+ // Leak the library when tracing, since the profiler may still be reading it.
+ // TODO(benvanik): move to an atexit handler instead, verify with ASAN/MSAN
+ executable_library_.release();
+#else
executable_library_.reset();
+#endif // IREE_TRACING_FEATURES & IREE_TRACING_FEATURE_INSTRUMENTATION
for (const auto& file_path : temp_file_paths_) {
file_io::DeleteFile(file_path).IgnoreError();
}
diff --git a/iree/test/e2e/regression/slice_add.mlir b/iree/test/e2e/regression/slice_add.mlir
new file mode 100644
index 0000000..762d979
--- /dev/null
+++ b/iree/test/e2e/regression/slice_add.mlir
@@ -0,0 +1,15 @@
+// RUN: iree-run-mlir -export-all -iree-hal-target-backends=vmla -function-input="3x4xi32=[[1,2,3,4],[5,6,7,8],[9,10,11,12]]" -function-input="1x2xi32=10" %s | IreeFileCheck %s
+// RUN: [[ $IREE_LLVMJIT_DISABLE == 1 ]] || (iree-run-mlir -export-all -iree-hal-target-backends=llvm-ir -function-input="3x4xi32=[[1,2,3,4],[5,6,7,8],[9,10,11,12]]" -function-input="1x2xi32=10" %s | IreeFileCheck %s)
+// RUN: [[ $IREE_VULKAN_DISABLE == 1 ]] || (iree-run-mlir -export-all -iree-hal-target-backends=vulkan-spirv -function-input="3x4xi32=[[1,2,3,4],[5,6,7,8],[9,10,11,12]]" -function-input="1x2xi32=10" %s | IreeFileCheck %s)
+
+// CHECK: EXEC @slice_stride_part
+// CHECK: 1x2xi32=[16 17]
+func @slice_stride_part(%arg0: tensor<3x4xi32>, %arg1: tensor<1x2xi32>) -> tensor<1x2xi32> {
+ %1 = "mhlo.slice"(%arg0) {
+ start_indices = dense<[1, 1]> : tensor<2xi64>,
+ limit_indices = dense<[2, 3]> : tensor<2xi64>,
+ strides = dense<1> : tensor<2xi64>
+ } : (tensor<3x4xi32>) -> tensor<1x2xi32>
+ %2 = mhlo.add %1, %arg1 : tensor<1x2xi32>
+ return %2 : tensor<1x2xi32>
+}
diff --git a/iree/vm/stack.h b/iree/vm/stack.h
index 6c50e7f..bf59de6 100644
--- a/iree/vm/stack.h
+++ b/iree/vm/stack.h
@@ -87,7 +87,7 @@
// code), etc.
iree_vm_source_offset_t pc;
- IREE_TRACE(iree_zone_id_t trace_zone)
+ IREE_TRACE(iree_zone_id_t trace_zone;)
} iree_vm_stack_frame_t;
// Returns the implementation-defined frame storage associated with |frame|.
diff --git a/scripts/update_e2e_coverage.py b/scripts/update_e2e_coverage.py
index 325cb28..fed8c12 100755
--- a/scripts/update_e2e_coverage.py
+++ b/scripts/update_e2e_coverage.py
@@ -28,13 +28,12 @@
TENSORFLOW_COVERAGE_DIR = 'tensorflow_coverage'
REFERENCE_BACKEND = 'tf'
-# Assumes that tests are expanded for the tf, iree_vmla, iree_llvmjit and
+# Assumes that tests are expanded for the tf, iree_vmla, and
# iree_vulkan backends.
BACKENDS_TO_TITLES = collections.OrderedDict([
('tf', 'tensorflow'),
('tflite', 'tflite'),
('iree_vmla', 'vmla'),
- ('iree_llvmjit', 'llvm-ir'),
('iree_vulkan', 'vulkan-spirv'),
])
