Merge pull request #3700 from hanhanW:main-to-google
PiperOrigin-RevId: 340471171
diff --git a/CMakeLists.txt b/CMakeLists.txt
index b2c30b7..bbff376 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -229,6 +229,14 @@
option(IREE_ENABLE_ASAN "Enable address sanitizer" OFF)
option(IREE_ENABLE_MSAN "Enable memory sanitizer" OFF)
option(IREE_ENABLE_TSAN "Enable thread sanitizer" OFF)
+option(IREE_ENABLE_CCACHE "Use ccache if installed to speed up rebuilds." OFF)
+
+if (${IREE_ENABLE_CCACHE})
+ find_program(CCACHE_PROGRAM ccache)
+ if(CCACHE_PROGRAM)
+ set_property(GLOBAL PROPERTY RULE_LAUNCH_COMPILE "${CCACHE_PROGRAM}")
+ endif()
+endif()
#-------------------------------------------------------------------------------
# IREE utility definitions
@@ -236,6 +244,7 @@
include(iree_macros)
include(iree_copts)
+include(sanitizers)
include(iree_whole_archive_link)
include(iree_cc_binary)
include(iree_cc_library)
@@ -261,7 +270,6 @@
# IREE compilation flags
#-------------------------------------------------------------------------------
-iree_append_list_to_string(CMAKE_CXX_FLAGS ${IREE_DEFAULT_COPTS})
iree_append_list_to_string(CMAKE_C_FLAGS_DEBUG ${IREE_C_FLAGS_DEBUG_LIST})
iree_append_list_to_string(CMAKE_CXX_FLAGS_DEBUG ${IREE_CXX_FLAGS_DEBUG_LIST})
diff --git a/build_tools/cmake/iree_copts.cmake b/build_tools/cmake/iree_copts.cmake
index 11c5f50..b6a8d76 100644
--- a/build_tools/cmake/iree_copts.cmake
+++ b/build_tools/cmake/iree_copts.cmake
@@ -135,54 +135,6 @@
set(IREE_TEST_COPTS "${ABSL_TEST_COPTS}")
#-------------------------------------------------------------------------------
-# Sanitizer configurations
-#-------------------------------------------------------------------------------
-
-include(CheckCXXCompilerFlag)
-if(${IREE_ENABLE_ASAN})
- set(CMAKE_REQUIRED_FLAGS "-fsanitize=address")
- check_cxx_compiler_flag(-fsanitize=address COMPILER_SUPPORTS_ASAN)
- unset(CMAKE_REQUIRED_FLAGS)
- if(${COMPILER_SUPPORTS_ASAN})
- list(APPEND IREE_DEFAULT_COPTS "-fsanitize=address")
- if(ANDROID)
- list(APPEND IREE_DEFAULT_LINKOPTS "-fuse-ld=gold")
- endif()
- else()
- message(FATAL_ERROR "The compiler does not support address sanitizer "
- "or is missing configuration for address sanitizer")
- endif()
-endif()
-if(${IREE_ENABLE_MSAN})
- set(CMAKE_REQUIRED_FLAGS "-fsanitize=memory")
- check_cxx_compiler_flag(-fsanitize=memory COMPILER_SUPPORTS_MSAN)
- unset(CMAKE_REQUIRED_FLAGS)
- if(${COMPILER_SUPPORTS_MSAN})
- list(APPEND IREE_DEFAULT_COPTS "-fsanitize=memory")
- if(ANDROID)
- list(APPEND IREE_DEFAULT_LINKOPTS "-fuse-ld=gold")
- endif()
- else()
- message(FATAL_ERROR "The compiler does not support memory sanitizer "
- "or is missing configuration for memory sanitizer")
- endif()
-endif()
-if(${IREE_ENABLE_TSAN})
- set(CMAKE_REQUIRED_FLAGS "-fsanitize=thread")
- check_cxx_compiler_flag(-fsanitize=thread COMPILER_SUPPORTS_TSAN)
- unset(CMAKE_REQUIRED_FLAGS)
- if(${COMPILER_SUPPORTS_TSAN})
- list(APPEND IREE_DEFAULT_COPTS "-fsanitize=thread")
- if(ANDROID)
- list(APPEND IREE_DEFAULT_LINKOPTS "-fuse-ld=gold")
- endif()
- else()
- message(FATAL_ERROR "The compiler does not support thread sanitizer "
- "or is missing configuration for thread sanitizer")
- endif()
-endif()
-
-#-------------------------------------------------------------------------------
# Size-optimized build flags
#-------------------------------------------------------------------------------
diff --git a/build_tools/cmake/iree_cross_compile.cmake b/build_tools/cmake/iree_cross_compile.cmake
index b8c1fde..3cfa1e9 100644
--- a/build_tools/cmake/iree_cross_compile.cmake
+++ b/build_tools/cmake/iree_cross_compile.cmake
@@ -124,6 +124,12 @@
# )
-DIREE_HAL_DRIVERS_TO_BUILD="${_CONFIG_HAL_DRIVERS_TO_BUILD}"
-DIREE_TARGET_BACKENDS_TO_BUILD="${_CONFIG_TARGET_BACKENDS_TO_BUILD}"
+ -DIREE_ENABLE_ASAN=${IREE_ENABLE_ASAN}
+ -DIREE_ENABLE_MSAN=${IREE_ENABLE_MSAN}
+ -DIREE_ENABLE_TSAN=${IREE_ENABLE_TSAN}
+ -DIREE_ENABLE_CCACHE=${IREE_ENABLE_CCACHE}
+ -DLLVM_CCACHE_BUILD=${LLVM_CCACHE_BUILD}
+ -DCMAKE_VERBOSE_MAKEFILE=${CMAKE_VERBOSE_MAKEFILE}
WORKING_DIRECTORY ${_CONFIG_BINARY_ROOT}
DEPENDS iree_prepare_${CONFIG_NAME}_dir
COMMENT "Configuring IREE for ${CONFIG_NAME} build...")
diff --git a/build_tools/cmake/sanitizers.cmake b/build_tools/cmake/sanitizers.cmake
new file mode 100644
index 0000000..026870c
--- /dev/null
+++ b/build_tools/cmake/sanitizers.cmake
@@ -0,0 +1,42 @@
+# 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.
+
+#-------------------------------------------------------------------------------
+# Sanitizer configurations
+#-------------------------------------------------------------------------------
+
+# Note: we add these flags to the global CMake flags, not to IREE-specific
+# variables such as IREE_DEFAULT_COPTS so that all symbols are consistently
+# defined with the same sanitizer flags, including e.g. standard library
+# symbols that might be used by both IREE and non-IREE (e.g. LLVM) code.
+
+if(${IREE_ENABLE_ASAN})
+ string(APPEND CMAKE_CXX_FLAGS " -fsanitize=address")
+ string(APPEND CMAKE_C_FLAGS " -fsanitize=address")
+endif()
+if(${IREE_ENABLE_MSAN})
+ string(APPEND CMAKE_CXX_FLAGS " -fsanitize=memory")
+ string(APPEND CMAKE_C_FLAGS " -fsanitize=memory")
+endif()
+if(${IREE_ENABLE_TSAN})
+ string(APPEND CMAKE_CXX_FLAGS " -fsanitize=thread")
+ string(APPEND CMAKE_C_FLAGS " -fsanitize=thread")
+endif()
+if(ANDROID)
+ # Work around https://github.com/android/ndk/issues/1088
+ if(${IREE_ENABLE_ASAN} OR ${IREE_ENABLE_MSAN} OR ${IREE_ENABLE_TSAN})
+ string(APPEND CMAKE_EXE_LINKER_FLAGS " -fuse-ld=gold")
+ string(APPEND CMAKE_SHARED_LINKER_FLAGS " -fuse-ld=gold")
+ endif()
+endif()
\ No newline at end of file
diff --git a/iree/compiler/Dialect/Flow/Transforms/CreateBenchmarkFuncs.cpp b/iree/compiler/Dialect/Flow/Transforms/CreateBenchmarkFuncs.cpp
index 836c4f2..cbfdebe 100644
--- a/iree/compiler/Dialect/Flow/Transforms/CreateBenchmarkFuncs.cpp
+++ b/iree/compiler/Dialect/Flow/Transforms/CreateBenchmarkFuncs.cpp
@@ -88,22 +88,31 @@
builder.getFunctionType({}, funcOp.getType().getResults());
std::string funcName = std::string(funcOp.getName()) + "_dummy_args";
auto newFuncOp = builder.create<FuncOp>(loc, funcName, funcType);
- newFuncOp.setAttr("iree.module.export", builder.getUnitAttr());
- Block* block = newFuncOp.addEntryBlock();
+ // Insert module-scope ops before the new function.
OpBuilder::InsertionGuard guard(builder);
builder.setInsertionPoint(newFuncOp);
- auto blockBuilder = OpBuilder::atBlockBegin(block);
- BlockAndValueMapping mapping;
- for (auto iter : llvm::enumerate(funcOp.getType().getInputs())) {
- auto arg = getDummyInput(builder, blockBuilder, loc, iter.value(),
- moduleSymbols);
- mapping.map(funcOp.getArgument(iter.index()), arg);
- }
- for (auto& op : funcOp.getRegion().begin()->getOperations()) {
- blockBuilder.clone(op, mapping);
- }
+ // Instead of operating on `newFuncOp`, we clone a FuncOp. This is because
+ // we need to set the mapper before cloning blocks. However, the mapping
+ // values must be created within a block, so it will be within the block
+ // scope. The implementation creates dummy inputs to the entry block of
+ // the cloned function, and replaces all uses with them. In the end, we
+ // erase all the arguments and clone the modified function into
+ // ` newFuncOp`.
+ FuncOp clonedFuncOp = funcOp.clone();
+ auto blockBuilder = OpBuilder::atBlockBegin(&(*clonedFuncOp.begin()));
+ for (int i = 0, e = clonedFuncOp.getNumArguments(); i < e; ++i) {
+ auto arg = clonedFuncOp.getArgument(i);
+ auto newArg = getDummyInput(builder, blockBuilder, loc, arg.getType(),
+ moduleSymbols);
+ arg.replaceAllUsesWith(newArg);
+ }
+ clonedFuncOp.eraseArguments(llvm::to_vector<4>(
+ llvm::seq<unsigned>(0, clonedFuncOp.getNumArguments())));
+ BlockAndValueMapping mapping;
+ clonedFuncOp.cloneInto(newFuncOp, mapping);
+ newFuncOp.setAttr("iree.module.export", builder.getUnitAttr());
funcOp.removeAttr("iree.module.export");
}
}
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/create_benchmark_funcs.mlir b/iree/compiler/Dialect/Flow/Transforms/test/create_benchmark_funcs.mlir
index b742302..f53085c 100644
--- a/iree/compiler/Dialect/Flow/Transforms/test/create_benchmark_funcs.mlir
+++ b/iree/compiler/Dialect/Flow/Transforms/test/create_benchmark_funcs.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -iree-flow-transformation-pipeline -iree-flow-export-dispatches %s | IreeFileCheck %s
+// RUN: iree-opt -split-input-file -iree-flow-transformation-pipeline -iree-flow-export-dispatches %s | IreeFileCheck %s
module {
func @two_dispatch(%arg0: tensor<5x3xf32>, %arg1: tensor<3x5xf32>) -> (tensor<5x5xf32>, tensor<3x5xf32>) attributes { iree.module.export } {
@@ -36,3 +36,23 @@
// CHECK: %[[DISPATCH_RES1:.+]] = flow.dispatch
// CHECK: %[[DISPATCH_RES2:.+]] = flow.dispatch
// CHECK: flow.return %[[DISPATCH_RES1]], %[[DISPATCH_RES2]] : tensor<5x5xf32>, tensor<3x5xf32>
+
+// -----
+
+func @while(%start: tensor<i32>, %bound: tensor<i32>) -> tensor<i32> attributes { iree.module.export } {
+ %res = "mhlo.while"(%start) ( {
+ ^bb0(%count: tensor<i32>):
+ %1 = "mhlo.compare"(%count, %bound) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
+ "mhlo.return"(%1) : (tensor<i1>) -> ()
+ }, {
+ ^bb0(%count: tensor<i32>):
+ %1 = mhlo.add %count, %count : tensor<i32>
+ "mhlo.return"(%1) : (tensor<i32>) -> ()
+ }) : (tensor<i32>) -> tensor<i32>
+ return %res : tensor<i32>
+}
+
+// CHECK: func @while_dummy_args
+// CHECK: %{{.+}} = flow.variable.load @{{.+}} : tensor<i32>
+// CHECK: %{{.+}} = flow.variable.load @{{.+}} : tensor<i32>
+// CHECK: br ^bb1
diff --git a/iree/hal/dylib/dylib_executable.cc b/iree/hal/dylib/dylib_executable.cc
index 4bf3c35..ef37edd 100644
--- a/iree/hal/dylib/dylib_executable.cc
+++ b/iree/hal/dylib/dylib_executable.cc
@@ -81,6 +81,9 @@
const auto& entry_points = *dylib_executable_def->entry_points();
entry_functions_.resize(entry_points.size());
+#if IREE_TRACING_FEATURES & IREE_TRACING_FEATURE_INSTRUMENTATION
+ entry_names_.resize(entry_points.size());
+#endif // IREE_TRACING_FEATURES & IREE_TRACING_FEATURE_INSTRUMENTATION
for (int i = 0; i < entry_functions_.size(); ++i) {
void* symbol = executable_library_->GetSymbol(entry_points[i]->c_str());
if (!symbol) {
diff --git a/scripts/android_symbolize.sh b/scripts/android_symbolize.sh
new file mode 100755
index 0000000..9dafab7
--- /dev/null
+++ b/scripts/android_symbolize.sh
@@ -0,0 +1,79 @@
+#!/bin/bash
+
+# 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.
+
+# Reads text from stdin containing stack frames in the format
+# `#3 0x5b09835a14 (/data/local/tmp/iree-benchmark-module+0x18aa14)`
+# and symbolizes such lines and prints them to stdout.
+#
+# The approach is to pull the files out of the device and use the host
+# build of llvm-symbolizer provided in the NDK.
+#
+# Other lines are just echo'd to stdout.
+# This is meant to be used like this:
+# ANDROID_NDK=~/android-ndk-r21d ~/android-symbolize.sh < ~/a
+# Where one has previously stored e.g. some ASan report to a file, here ~/a.
+#
+# Discussion of alternatives: https://github.com/android/ndk/issues/753
+
+# Provide the location of your Android NDK in the ANDROID_NDK env var.
+: ${ANDROID_NDK:=""}
+
+if [ -z "${ANDROID_NDK}" ]
+then
+ echo "Please define ANDROID_NDK to point to your Android NDK."
+ exit 1
+fi
+
+tmpdir="$(mktemp -d)"
+
+while read line
+do
+ header="$(echo "$line" | grep -o '^\s*#[0-9]\+')"
+ if [ -z "$header" ]
+ then
+ echo "$line"
+ continue
+ fi
+
+ location_with_parentheses="$(echo "$line" | grep -o '(/[^)]*)$')"
+ location="$(echo "$location_with_parentheses" | tr -d '()' )"
+ file="$(echo "$location" | cut -d '+' -f 1)"
+ address="$(echo "$location" | cut -d '+' -f 2)"
+ if [[ -z "$file" || -z "$address" ]]
+ then
+ echo "$line"
+ continue
+ fi
+
+ file_basename="$(basename "$file")"
+ pulled_file="$tmpdir/$file_basename"
+ if [ ! -f "$pulled_file" ]
+ then
+ adb pull "$file" "$pulled_file" 1>/dev/null 2>/dev/null
+ fi
+ llvm_symbolizer_output="$($ANDROID_NDK/toolchains/llvm/prebuilt/linux-x86_64/bin/llvm-symbolizer -obj "$pulled_file" "$address")"
+ if [ -z "$llvm_symbolizer_output" ]
+ then
+ echo "$line"
+ continue
+ fi
+
+ function="$(echo "$llvm_symbolizer_output" | head -n1)"
+ source_location="$(echo "$llvm_symbolizer_output" | tail -n1)"
+ echo "$header $source_location $function"
+done
+
+rm -rf "$tmpdir"
diff --git a/scripts/android_tracy_tutorial.sh b/scripts/android_tracy_tutorial.sh
new file mode 100755
index 0000000..812c09e
--- /dev/null
+++ b/scripts/android_tracy_tutorial.sh
@@ -0,0 +1,261 @@
+#!/bin/bash
+
+# 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.
+
+# This script assumes Debian Linux. It could easily be modified to run on any Unix,
+# as the Debian assumption is only to automatically install some packages,
+# (Tracy dependencies). With a bit more work it could be made to support other OSes
+# such as Windows, but that's not at all handled yet.
+
+function print_status {
+ echo -e "\e[96m$@\e[39m"
+}
+
+# Environment variables. They can be set manually, or will use the following defaults.
+# IREE_ROOT and ANDROID_NDK are empty by default, must be defined by the user.
+#
+# To make this script very easy to play with, this script will also default to
+# benchmarking a specific NN model (currently MobileBert encoder), and will take
+# care of generating its input MLIR form and of setting the right --function_inputs
+# and --entry_function for it. To use this script on any other specific NN model,
+# define the following environment variables:
+# IREE_INPUT_MLIR
+# FUNCTION_INPUTS
+# ENTRY_FUNCTION
+
+DEFAULT_IREE_INPUT_MLIR="/tmp/iree/modules/MobileBertSquad/iree_input.mlir"
+
+: ${IREE_ROOT:=""}
+: ${IREE_BUILD_ANDROID:="$HOME/iree-build-android"}
+: ${TRACY_ROOT:="$HOME/tracy"}
+: ${PYTHON_BIN:=python3}
+: ${CC:=clang}
+: ${CXX:=clang++}
+: ${IREE_INPUT_MLIR:="${DEFAULT_IREE_INPUT_MLIR}"}
+: ${ANDROID_NDK:=""}
+: ${IREE_LLVMAOT_LINKER_PATH:="$ANDROID_NDK/toolchains/llvm/prebuilt/linux-x86_64/bin/aarch64-linux-android30-clang++ -static-libstdc++ -O3"}
+: ${FUNCTION_INPUTS:=""}
+: ${ENTRY_FUNCTION:=""}
+# Validation of the environment variables.
+
+if [ -z "${IREE_ROOT}" ]
+then
+ print_status "Please define IREE_ROOT to point to your IREE git clone."
+ print_status "Example:"
+ print_status " IREE_ROOT=\$HOME/iree $0"
+ exit 1
+fi
+
+if [ ! -f "${IREE_ROOT}/iree/base/CMakeLists.txt" ]
+then
+ print_status "FATAL: bad IREE_ROOT (${IREE_ROOT})."
+ exit 1
+fi
+
+if [ -z "${ANDROID_NDK}" ]
+then
+ print_status "Please define ANDROID_NDK to point to your Android NDK."
+ print_status "Example:"
+ print_status " ANDROID_NDK=\$HOME/android-ndk-r21d $0"
+ exit 1
+fi
+
+if [ ! -d "${ANDROID_NDK}/toolchains/llvm" ]
+then
+ print_status "FATAL: bad ANDROID_NDK (${ANDROID_NDK})."
+ exit 1
+fi
+
+if ! $CC --version 1>/dev/null 2>/dev/null
+then
+ print_status "FATAL: Install $CC and set CC to it."
+ exit 1
+fi
+
+if ! $CXX --version 1>/dev/null 2>/dev/null
+then
+ print_status "FATAL: Install $CXX and set CXX to it."
+ exit 1
+fi
+
+if ! $IREE_LLVMAOT_LINKER_PATH --version 1>/dev/null 2>/dev/null
+then
+ print_status "Bad IREE_LLVMAOT_LINKER_PATH (${IREE_LLVMAOT_LINKER_PATH}). Rerun with it correctly set."
+ exit 1
+fi
+
+# If we're playing with the default input MLIR, preset some reasonable
+# FUNCTION_INPUTS and ENTRY_FUNCTION if they haven't been set.
+# Otherwise, they must be provided.
+if [[ -z "${FUNCTION_INPUTS}" || -z "${ENTRY_FUNCTION}" ]]
+then
+ if [[ "${IREE_INPUT_MLIR}" == "${DEFAULT_IREE_INPUT_MLIR}" ]]
+ then
+ FUNCTION_INPUTS="1x384xi32,1x384xi32,1x384xi32"
+ ENTRY_FUNCTION="serving_default"
+ else
+ print_status "Please specify FUNCTION_INPUTS and ENTRY_FUNCTION appropriately for your IREE_INPUT_MLIR (${IREE_INPUT_MLIR})"
+ fi
+fi
+
+print_status "Running with the following environment variables:"
+print_status "ANDROID_NDK=${ANDROID_NDK}"
+print_status "IREE_LLVMAOT_LINKER_PATH=${IREE_LLVMAOT_LINKER_PATH}"
+print_status "IREE_ROOT=${IREE_ROOT}"
+print_status "IREE_BUILD_ANDROID=${IREE_BUILD_ANDROID}"
+print_status "TRACY_ROOT=${TRACY_ROOT}"
+print_status "PYTHON_BIN=${PYTHON_BIN}"
+print_status "CC=${CC}"
+print_status "CXX=${CXX}"
+print_status "IREE_INPUT_MLIR=${IREE_INPUT_MLIR}"
+print_status "FUNCTION_INPUTS=${FUNCTION_INPUTS}"
+print_status "ENTRY_FUNCTION=${ENTRY_FUNCTION}"
+
+echo
+
+print_status "Ensuring that we have Tracy source code..."
+if [ ! -d ${TRACY_ROOT}/profiler/build/unix/ ]
+then
+ print_status "Tracy not found at ${TRACY_ROOT}. Either set TRACY_ROOT, or we're going to clone the git repository now."
+ read -p "Press the return key..."
+ git clone https://github.com/wolfpld/tracy "${TRACY_ROOT}"
+fi
+
+echo
+
+print_status "Ensuring that the Tracy profiler is built..."
+if [ ! -x "${TRACY_ROOT}/profiler/build/unix/Tracy-release" ]
+then
+ print_status "Checking Tracy dependencies - assuming Debian."
+ TRACY_DEPS="libcapstone-dev libtbb-dev libglfw3-dev libfreetype6-dev libgtk-3-dev"
+ TRACY_DEPS_COUNT=5
+ TRACY_DEPS_INSTALLED="$(apt list $TRACY_DEPS) 2>/dev/null | grep installed | wc -l"
+ if [ $TRACY_DEPS_INSTALLED != $TRACY_DEPS_COUNT ]
+ then
+ print_status "Installing dependencies now - assuming Debian."
+ sudo apt install $TRACY_DEPS
+ fi
+ make -C "${TRACY_ROOT}/profiler/build/unix" -j12 release
+fi
+
+echo
+
+print_status "Ensuring that we have the input MLIR file..."
+if [ ! -f "${IREE_INPUT_MLIR}" ]
+then
+ print_status "Set IREE_INPUT_MLIR to point to some iree input MLIR file. Not found at current value ${IREE_INPUT_MLIR}."
+ if [ "${IREE_INPUT_MLIR}" == "${DEFAULT_IREE_INPUT_MLIR}" ]
+ then
+ print_status "Okay, we actually know how to generate that file, ${IREE_INPUT_MLIR}, but it will take a while."
+ print_status "Press the return key to continue..."
+ pushd "${IREE_ROOT}"
+ scripts/get_e2e_artifacts.py --test_suites=mobile_bert_squad_tests
+ popd
+ fi
+fi
+
+if [ ! -f "${IREE_INPUT_MLIR}" ]
+then
+ print_status "FATAL: we should have ${IREE_INPUT_MLIR} by that point."
+ exit 1
+fi
+
+echo
+
+print_status "Building IREE for Android in ${IREE_BUILD_ANDROID}..."
+
+mkdir -p "${IREE_BUILD_ANDROID}"
+pushd "${IREE_BUILD_ANDROID}"
+
+cmake -G Ninja ../iree \
+ -DCMAKE_TOOLCHAIN_FILE=$ANDROID_NDK/build/cmake/android.toolchain.cmake \
+ -DANDROID_ABI="arm64-v8a" \
+ -DANDROID_PLATFORM=android-30 \
+ -DCMAKE_BUILD_TYPE=RelWithDebInfo \
+ -DIREE_BUILD_COMPILER=OFF \
+ -DIREE_BUILD_SAMPLES=OFF \
+ -DIREE_HOST_C_COMPILER=`which "$CC"` \
+ -DIREE_HOST_CXX_COMPILER=`which "$CXX"` \
+ -DIREE_ENABLE_RUNTIME_TRACING=ON
+
+cmake --build . --target \
+ iree_tools_iree-translate \
+ iree_tools_iree-benchmark-module
+
+popd
+echo
+
+print_status "Compiling the input MLIR file into a IREE module..."
+
+IREE_COMPILED_MODULE=/tmp/android_module.vmfb
+IREE_LOG=/tmp/iree-translate.log
+
+rm -rf "${IREE_COMPILED_MODULE}"
+
+IREE_LLVMAOT_LINKER_PATH="${IREE_LLVMAOT_LINKER_PATH}" \
+ "${IREE_BUILD_ANDROID}/host/bin/iree-translate" \
+ --iree-hal-target-backends=dylib-llvm-aot \
+ --iree-mlir-to-vm-bytecode-module \
+ --iree-llvm-target-triple=aarch64-linux-android \
+ /tmp/iree/modules/MobileBertSquad/iree_input.mlir \
+ -o "${IREE_COMPILED_MODULE}" \
+ 2>"${IREE_LOG}"
+
+if [ ! -f "${IREE_COMPILED_MODULE}" ]
+then
+ print_status "iree-translate failed to produce ${IREE_COMPILED_MODULE}. Log saved in ${IREE_LOG}. First few lines:"
+ # The whole log might be enormous if it contains a big MLIR dump.
+ head -n10 "${IREE_LOG}"
+ print_status "tip: check if IREE_LLVMAOT_LINKER_PATH was correctly set."
+ exit 1
+fi
+
+echo
+
+DEVICE_IREE_COMPILED_MODULE=/data/local/tmp/android_module.vmfb
+
+print_status "Pushing the compiled module to the device..."
+adb push "${IREE_COMPILED_MODULE}" "${DEVICE_IREE_COMPILED_MODULE}"
+echo
+
+print_status "Pushing the IREE benchmarking program to the device..."
+adb push "${IREE_BUILD_ANDROID}"/iree/tools/iree-benchmark-module /data/local/tmp
+echo
+
+print_status "Setting up TCP port forwarding to let Tracy connect with the benchmark running on the device..."
+adb forward tcp:8086 tcp:8086
+echo
+
+print_status "Now you can launch the Tracy UI in another shell and hit \"Connect\", while we run the benchmark on the device in this shell..."
+print_status "Run this command in another shell:"
+print_status " ${TRACY_ROOT}/profiler/build/unix/Tracy-release"
+echo
+
+print_status "Running the benchmark... hit Ctrl-C to terminate it after Tracy is done with it."
+# `TRACY_NO_EXIT=1` is to prevent it from exiting at the end: that's needed for profiling
+# short-running tasks.
+# `taskset 80` selects which CPU core to run on. On Pixel4, `taskset 80` gives the biggest
+# core, which can get some reproducibility, as long as we don't run into thermal issues.
+# `taskset 0f` would give the little cores, avoiding thermal issues but running slower and
+# requiring different optimization work to be efficient on.
+# `--driver=dylib` is to use the LLVM AOT generated code backend.
+adb shell \
+ TRACY_NO_EXIT=1 \
+ taskset 80 \
+ data/local/tmp/iree-benchmark-module \
+ --driver=dylib \
+ --module_file="${DEVICE_IREE_COMPILED_MODULE}" \
+ --function_inputs="${FUNCTION_INPUTS}" \
+ --entry_function="${ENTRY_FUNCTION}"
