From 1c5d557a36ba66f00574c4141cadc6ed6180e3f4 Mon Sep 17 00:00:00 2001 From: Scott Todd Date: Wed, 1 Nov 2023 11:29:05 -0700 Subject: [PATCH 01/16] Add new docs/developers/ section to website (empty). --- docs/website/docs/developers/index.md | 3 +++ docs/website/mkdocs.yml | 2 ++ 2 files changed, 5 insertions(+) create mode 100644 docs/website/docs/developers/index.md diff --git a/docs/website/docs/developers/index.md b/docs/website/docs/developers/index.md new file mode 100644 index 000000000000..aedbb284c9c8 --- /dev/null +++ b/docs/website/docs/developers/index.md @@ -0,0 +1,3 @@ +# Developers + +TODO diff --git a/docs/website/mkdocs.yml b/docs/website/mkdocs.yml index 4867083a15ae..6249c23d97c1 100644 --- a/docs/website/mkdocs.yml +++ b/docs/website/mkdocs.yml @@ -165,6 +165,8 @@ nav: - Glossary: "reference/glossary.md" - Optimization options: "reference/optimization-options.md" - Extensions: "reference/extensions.md" + - "Developers": + - "developers/index.md" - "Community": - "community/index.md" - "Blog": From 7607d867da22f0b875ee3d5c4c7d28c6b4e67432 Mon Sep 17 00:00:00 2001 From: Scott Todd Date: Wed, 1 Nov 2023 11:29:40 -0700 Subject: [PATCH 02/16] Move "get_started" docs to website/docs/developers. --- build_tools/scripts/integrate/README.md | 4 +- docs/developers/get_started/README.md | 12 - .../get_started/building_with_bazel_linux.md | 117 ---------- .../get_started/building_with_bazel_macos.md | 120 ---------- .../building_with_bazel_windows.md | 120 ---------- .../docs/developers/building_with_bazel.md | 219 ++++++++++++++++++ .../developers}/building_with_emscripten.md | 10 +- .../cmake_options_and_variables.md | 102 +++++--- .../developers}/vulkan_environment_setup.md | 2 +- docs/website/mkdocs.yml | 4 + 10 files changed, 307 insertions(+), 403 deletions(-) delete mode 100644 docs/developers/get_started/README.md delete mode 100644 docs/developers/get_started/building_with_bazel_linux.md delete mode 100644 docs/developers/get_started/building_with_bazel_macos.md delete mode 100644 docs/developers/get_started/building_with_bazel_windows.md create mode 100644 docs/website/docs/developers/building_with_bazel.md rename docs/{developers/get_started => website/docs/developers}/building_with_emscripten.md (91%) rename docs/{developers/get_started => website/docs/developers}/cmake_options_and_variables.md (71%) rename docs/{developers/get_started => website/docs/developers}/vulkan_environment_setup.md (99%) diff --git a/build_tools/scripts/integrate/README.md b/build_tools/scripts/integrate/README.md index a83799e54626..f42de54d03c1 100644 --- a/build_tools/scripts/integrate/README.md +++ b/build_tools/scripts/integrate/README.md @@ -290,7 +290,9 @@ cmake -G Ninja \ .. ``` -To repro failures in CI `bazel_linux_x86-swiftshader_core`, we can follow the [doc](https://github.com/openxla/iree/blob/main/docs/developers/get_started/building_with_bazel_linux.md) to build IREE using bazel. E.g., +To repro failures in CI `bazel_linux_x86-swiftshader_core`, we can follow the +[developer doc](https://iree.dev/developers/bazel) to build IREE using bazel. +E.g., ```bash export CC=clang diff --git a/docs/developers/get_started/README.md b/docs/developers/get_started/README.md deleted file mode 100644 index 74b2b91b73bb..000000000000 --- a/docs/developers/get_started/README.md +++ /dev/null @@ -1,12 +0,0 @@ -# Additional getting started guides - ---- - -The primary guides are located at - (source in -[the website/ folder](../../website/docs/building-from-source/) ) - ---- - -The files in this folder contain a mix of legacy and minimal-effort (Bazel) -documentation. diff --git a/docs/developers/get_started/building_with_bazel_linux.md b/docs/developers/get_started/building_with_bazel_linux.md deleted file mode 100644 index 27126b3b9e55..000000000000 --- a/docs/developers/get_started/building_with_bazel_linux.md +++ /dev/null @@ -1,117 +0,0 @@ -# Getting Started on Linux with Bazel - -**NOTE** Bazel build support is primarily for internal project infrastructure. -We strongly recommend users build with CMake instead. - -This guide walks through building the core compiler and runtime parts of IREE -from source. Auxiliary components like the Python bindings and Vulkan driver are -documented separately, as they require further setup. - -## Prerequisites - -### Install Bazel - -Install Bazel, matching IREE's -[`.bazelversion`](https://github.com/openxla/iree/blob/main/.bazelversion) by -following the -[official docs](https://docs.bazel.build/versions/master/install.html). - -### Install a Compiler - -We recommend Clang. GCC is not fully supported. - -```shell -sudo apt install clang -``` - -Set environment variables for Bazel: - -```shell -export CC=clang -export CXX=clang++ -``` - -### Install python3 numpy - -```shell -python3 -m pip install numpy -``` - -## Clone and Build - -### Clone - -Clone the repository, initialize its submodules and configure: - -```shell -git clone https://github.com/openxla/iree.git -cd iree -git submodule update --init -python3 configure_bazel.py -``` - -> Tip:
->     Editors and other programs can also clone the -> repository, just make sure that they initialize the submodules. - -### Build - -Run all core tests: - -```shell -bazel test -k //... -``` - -> Tip:
->     You can add flags like -> `--test_env=IREE_VULKAN_DISABLE=1` to your test command to change how/which -> tests run. - -In general, build artifacts will be under the `bazel-bin` directory at the top -level. - -## Recommended user.bazelrc - -You can put a user.bazelrc at the root of the repository and it will be ignored -by git. The recommended contents for Linux are: - -```shell -build --disk_cache=/tmp/bazel-cache - -# Use --config=debug to compile IREE and LLVM without optimizations -# and with assertions enabled. -build:debug --config=asserts --compilation_mode=opt '--per_file_copt=iree|llvm@-O0' --strip=never - -# Use --config=asserts to enable assertions. This has to be done globally: -# Code compiled with and without assertions can't be linked together (ODR violation). -build:asserts --compilation_mode=opt '--copt=-UNDEBUG' -``` - -## What's next? - -### Take a Look Around - -Build all of IREE's 'tools' directory: - -```shell -bazel build tools/... -``` - -Check out what was built: - -```shell -ls bazel-bin/tools/ -./bazel-bin/tools/iree-compile --help -``` - -Translate a -[MLIR file](https://github.com/openxla/iree/blob/main/samples/models/simple_abs.mlir) -and execute a function in the compiled module: - -```shell -# iree-run-mlir [input.mlir] -$ ./bazel-bin/tools/iree-run-mlir \ - --iree-hal-target-backends=vmvx --print-mlir \ - ./samples/models/simple_abs.mlir \ - --input=f32=-2 -``` diff --git a/docs/developers/get_started/building_with_bazel_macos.md b/docs/developers/get_started/building_with_bazel_macos.md deleted file mode 100644 index 884aa20d58a1..000000000000 --- a/docs/developers/get_started/building_with_bazel_macos.md +++ /dev/null @@ -1,120 +0,0 @@ -# Getting Started on macOS with Bazel - -**NOTE** Bazel build support is primarily for internal project infrastructure. -We strongly recommend users build with CMake instead. - -This guide walks through building the core compiler and runtime parts of IREE -from source. Auxiliary components like the Python bindings and Vulkan driver are -not documented for macOS at this time. - -IREE is not officially supported on macOS at this time. It may work, but it is -not a part of our open source CI, and may be intermittently broken. -Contributions related to macOS support and documentation are welcome however. - -## Prerequisites - -### Install Homebrew - -This guide uses [Homebrew](https://brew.sh/) to install IREE's dependencies. - -```shell -/bin/bash -c "$(curl -fsSL https://raw.githubusercontent.com/Homebrew/install/master/install.sh)" -``` - -### Install Bazel - -Install Bazel via Homebrew: - -```shell -brew install bazel -``` - -Note: when you first run `bazel` to build IREE, it will prompt you to copy and -run a shell command to select the right version. - -### Install python3 numpy - -```shell -python3 -m pip install numpy --user -``` - -## Clone and Build - -### Clone - -Clone the repository, initialize its submodules and configure: - -```shell -git clone https://github.com/openxla/iree.git -cd iree -git submodule update --init -python3 configure_bazel.py -``` - -> Tip:
->     Editors and other programs can also clone the -> repository, just make sure that they initialize the submodules. - -### Build - -Run all core tests that pass on our OSS CI: - -```shell -$ bazel test -k //... \ - --test_env=IREE_VULKAN_DISABLE=1 \ - --build_tag_filters="-nokokoro" \ - --test_tag_filters="--nokokoro,-driver=vulkan" -``` - -> Tip:
->     Not all tests are passing on macOS, but the build does -> complete successfully at the time of writing. - -In general, build artifacts will be under the `bazel-bin` directory at the top -level. - -## Recommended user.bazelrc - -You can put a user.bazelrc at the root of the repository and it will be ignored -by git. The recommended contents for Linux/macOS are: - -```shell -build --disk_cache=/tmp/bazel-cache - -# Use --config=debug to compile IREE and LLVM without optimizations -# and with assertions enabled. -build:debug --config=asserts --compilation_mode=opt '--per_file_copt=iree|llvm@-O0' --strip=never - -# Use --config=asserts to enable assertions. This has to be done globally: -# Code compiled with and without assertions can't be linked together (ODR violation). -build:asserts --compilation_mode=opt '--copt=-UNDEBUG' -``` - -## What's next? - -### Take a Look Around - -Build all of IREE's 'tools' directory: - -```shell -bazel build tools/... -``` - -Check out what was built: - -```shell -ls bazel-bin/tools/ -./bazel-bin/tools/iree-compile --help -``` - -Translate a -[MLIR file](https://github.com/openxla/iree/blob/main/samples/models/simple_abs.mlir) -and execute a function in the compiled module: - -```shell -# iree-run-mlir [input.mlir] -$ ./bazel-bin/tools/iree-run-mlir \ - --iree-hal-target-backends=vmvx --print-mlir \ - ./samples/models/simple_abs.mlir \ - --input=f32=-2 -``` diff --git a/docs/developers/get_started/building_with_bazel_windows.md b/docs/developers/get_started/building_with_bazel_windows.md deleted file mode 100644 index b67cafdc41ee..000000000000 --- a/docs/developers/get_started/building_with_bazel_windows.md +++ /dev/null @@ -1,120 +0,0 @@ -# Getting Started on Windows with Bazel - -**NOTE** Bazel build support is primarily for internal project infrastructure. -Bazel on Windows in particular is particularly unstable and unsupported. -We strongly recommend users build with CMake instead. - -This guide walks through building the core compiler and runtime parts of IREE -from source. Auxiliary components like the Python bindings and Vulkan driver are -documented separately, as they require further setup. - -## Prerequisites - -> Tip:
->     You can simplify installation by using a package -> manager like [Scoop](https://scoop.sh/) or -> [Chocolatey](https://chocolatey.org/). - -### Install Bazel - -Install Bazel version > 2.0.0 (see -[`.bazelversion`](https://github.com/openxla/iree/blob/main/.bazelversion) for -the specific version IREE uses) by following the -[official docs](https://docs.bazel.build/versions/master/install-windows.html). - -Also install [MSYS2](https://www.msys2.org/) by following Bazel's documentation. - -### Install Python3 - -Instructions for installation can be found -[here](https://www.python.org/downloads/windows/). - -### Install Build Tools For Visual Studio - -Install the full Visual Studio or "Build Tools For Visual Studio" from the -[downloads page](https://visualstudio.microsoft.com/downloads/). - -Set a few environment variables. You are welcome to configure these however you -choose. For example, you could set them as system or user level environment -variables through your "System Properties" or you could use a shell such as -PowerShell or [cmder](https://cmder.net/)). Setting them through PowerShell -would look like this: - -```powershell -> $env:BAZEL_VS = "C:\Program Files (x86)\Microsoft Visual Studio\2019\BuildTools" -``` - -## Clone and Build - -### Clone - -Using your shell of choice (such as PowerShell or [cmder](https://cmder.net/)), -clone the repository, initialize its submodules, and configure: - -```powershell -> git clone https://github.com/openxla/iree.git -> cd iree -> git submodule update --init -> python configure_bazel.py -``` - -> Tip:
->     Clone to a short path like `C:\projects\` to avoid -> issues with Windows maximum path lengths (260 characters). - -> Tip:
->     Editors and other programs can also clone the -> repository, just make sure that they initialize the submodules. - -> Tip:
->     configure_bazel.py only detects that you have Windows -> and will output the default `--config=windows` to `configured.bazelrc`, which -> assumes the latest version of MSVC. To avoid some warnings, you may want to -> replace it with `--config=msvc2017`. - -### Build - -Run all core tests: - -```powershell -> bazel test -k //... -``` - -In general, build artifacts will be under the `bazel-bin` directory at the top -level. - -## Recommended user.bazelrc - -You can put a user.bazelrc at the root of the repository and it will be ignored -by git. The recommended contents for Windows are: - -```starlark -build --disk_cache=c:/bazelcache -build:debug --compilation_mode=dbg --copt=/O2 --per_file_copt=iree@/Od --strip=never -``` - -## What's next? - -### Take a Look Around - -Build all of IREE's 'tools' directory: - -```powershell -> bazel build tools/... -``` - -Check out what was built: - -```powershell -> dir bazel-bin\iree\tools\ -> .\bazel-bin\tools\iree-compile.exe --help -``` - -Translate a -[MLIR file](https://github.com/openxla/iree/blob/main/samples/models/simple_abs.mlir) -and execute a function in the compiled module: - -```powershell -> REM iree-run-mlir [input.mlir] -> .\bazel-bin\tools\iree-run-mlir.exe --iree-hal-target-backends=vmvx --print-mlir .\iree\samples\models\simple_abs.mlir --input=f32=-2 -``` diff --git a/docs/website/docs/developers/building_with_bazel.md b/docs/website/docs/developers/building_with_bazel.md new file mode 100644 index 000000000000..d2e9f8efd0d1 --- /dev/null +++ b/docs/website/docs/developers/building_with_bazel.md @@ -0,0 +1,219 @@ +# Building with Bazel + +This page walks through building IREE from source using the +[Bazel build system](https://bazel.build/). + +!!! warning + + Bazel build support is primarily for internal project infrastructure. We + strongly recommend [using CMake](../building-from-source/index.md) instead. + + Our Bazel configuration is also _only_ tested on Linux. Windows and macOS + may be unstable. + +## :octicons-download-16: Prerequisites + +=== ":fontawesome-brands-linux: Linux" + + 1. Install Bazel, matching IREE's + [`.bazelversion`](https://github.com/openxla/iree/blob/main/.bazelversion) + by following the + [official docs](https://bazel.build/install). + + 2. Install a compiler such as Clang (GCC is not fully supported). + + ```shell + sudo apt install clang + ``` + + Set environment variables for Bazel: + + ```shell + export CC=clang + export CXX=clang++ + ``` + + 3. Install Python build requirements: + + ```shell + python -m pip install -r runtime/bindings/python/iree/runtime/build_requirements.txt + ``` + +=== ":fontawesome-brands-apple: macOS" + + 1. Install [Homebrew](https://brew.sh/): + + ```shell + /bin/bash -c "$(curl -fsSL https://raw.githubusercontent.com/Homebrew/install/master/install.sh)" + ``` + + 2. Install Bazel, matching IREE's + [`.bazelversion`](https://github.com/openxla/iree/blob/main/.bazelversion) + by following the [official docs](https://bazel.build/install/os-x) or + via Homebrew: + + ```shell + brew install bazel + ``` + + 3. Install Python build requirements: + + ```shell + python -m pip install -r runtime/bindings/python/iree/runtime/build_requirements.txt + ``` + +=== ":fontawesome-brands-windows: Windows" + + !!! tip + + You can simplify installation by using a package manager like + [Scoop](https://scoop.sh/) or [Chocolatey](https://chocolatey.org/). + + 1. Install Bazel, matching IREE's + [`.bazelversion`](https://github.com/openxla/iree/blob/main/.bazelversion) + by following the [official docs](https://bazel.build/install/windows). + + Also install [MSYS2](https://www.msys2.org/) by following Bazel's documentation. + + 2. Install Python3 ([docs here](https://www.python.org/downloads/windows/)) + and Python build requirements: + + ```shell + python -m pip install -r runtime/bindings/python/iree/runtime/build_requirements.txt + ``` + + 3. Install the full Visual Studio or "Build Tools For Visual Studio" from the + [downloads page](https://visualstudio.microsoft.com/downloads/) then + set the `BAZEL_VS` environment variable: + + ```powershell + > $env:BAZEL_VS = "C:\Program Files (x86)\Microsoft Visual Studio\2022\BuildTools" + ``` + +## :octicons-rocket-16: Quickstart: clone and build + +### Clone + +Use [Git](https://git-scm.com/) to clone the IREE repository and initialize its +submodules: + +```shell +git clone https://github.com/openxla/iree.git +cd iree +git submodule update --init +``` + +Configure Bazel: + +```shell +# This generates a `configured.bazelrc` file by analyzing your environment. +# Skipping this step will make it difficult to select your platform/compiler. +python3 configure_bazel.py +``` + +=== ":fontawesome-brands-linux: Linux" + + (No Linux-specific tips for configuring) + +=== ":fontawesome-brands-apple: macOS" + + (No macOS-specific tips for configuring) + +=== ":fontawesome-brands-windows: Windows" + + !!! tip + + Clone to a short path like `C:\projects\` to avoid issues with Windows + maximum path lengths (260 characters). + + !!! tip + + `configure_bazel.py` only detects that you have Windows and will output + the default `--config=windows` to `configured.bazelrc`, which assumes + the latest version of MSVC. To avoid some warnings, you may want to + replace it with (for example) `--config=msvc2022`. + +### Build + +Run all core tests: + +```shell +bazel test -k //... +``` + +!!! tip + + You can add flags like `--test_env=IREE_VULKAN_DISABLE=1` to your test + command to change how/which tests run. + +In general, build artifacts will be under the `bazel-bin` directory at the top +level. + +## :octicons-gear-16: Recommended `user.bazelrc` + +You can put a user.bazelrc at the root of the repository and it will be ignored +by git. + +=== ":fontawesome-brands-linux: Linux" + + ```shell + build --disk_cache=/tmp/bazel-cache + + # Use --config=debug to compile IREE and LLVM without optimizations + # and with assertions enabled. + build:debug --config=asserts --compilation_mode=opt '--per_file_copt=iree|llvm@-O0' --strip=never + + # Use --config=asserts to enable assertions. This has to be done globally: + # Code compiled with and without assertions can't be linked together (ODR violation). + build:asserts --compilation_mode=opt '--copt=-UNDEBUG' + ``` + +=== ":fontawesome-brands-apple: macOS" + + ```shell + build --disk_cache=/tmp/bazel-cache + + # Use --config=debug to compile IREE and LLVM without optimizations + # and with assertions enabled. + build:debug --config=asserts --compilation_mode=opt '--per_file_copt=iree|llvm@-O0' --strip=never + + # Use --config=asserts to enable assertions. This has to be done globally: + # Code compiled with and without assertions can't be linked together (ODR violation). + build:asserts --compilation_mode=opt '--copt=-UNDEBUG' + ``` + +=== ":fontawesome-brands-windows: Windows" + + ```shell + build --disk_cache=c:/bazelcache + build:debug --compilation_mode=dbg --copt=/O2 --per_file_copt=iree@/Od --strip=never + ``` + +## What's next? + +### Take a Look Around + +Build all of IREE's 'tools' directory: + +```shell +bazel build tools/... +``` + +Check out what was built: + +```shell +ls bazel-bin/tools/ +./bazel-bin/tools/iree-compile --help +``` + +Translate a +[MLIR file](https://github.com/openxla/iree/blob/main/samples/models/simple_abs.mlir) +and execute a function in the compiled module: + +```shell +# iree-run-mlir [input.mlir] +$ ./bazel-bin/tools/iree-run-mlir \ + --iree-hal-target-backends=vmvx --print-mlir \ + ./samples/models/simple_abs.mlir \ + --input=f32=-2 +``` diff --git a/docs/developers/get_started/building_with_emscripten.md b/docs/website/docs/developers/building_with_emscripten.md similarity index 91% rename from docs/developers/get_started/building_with_emscripten.md rename to docs/website/docs/developers/building_with_emscripten.md index 181a27120d6b..23238597a4bc 100644 --- a/docs/developers/get_started/building_with_emscripten.md +++ b/docs/website/docs/developers/building_with_emscripten.md @@ -1,4 +1,4 @@ -# Getting Started With Emscripten +# Building with Emscripten [Emscripten](https://emscripten.org/index.html) is a complete compiler toolchain to WebAssembly, using LLVM, with a special focus on speed, size, and @@ -24,9 +24,9 @@ Read and run source ./emsdk_env.sh ``` -## Building IREE's Runtime with Emscripten +## Building IREE's runtime with Emscripten -### Host Configuration +### Host configuration Build and install at least the compiler tools on your host machine, or install them from a binary distribution: @@ -40,7 +40,7 @@ $ cmake -G Ninja -B ../iree-build-host/ \ $ cmake --build ../iree-build-host/ --target install ``` -### Target Configuration +### Target configuration ```shell $ emcmake cmake -G Ninja -B ../iree-build-emscripten/ \ @@ -58,7 +58,7 @@ cmake --build ../iree-build-emscripten/ \ --target iree_samples_simple_embedding_simple_embedding_vmvx_sync ``` -### Load into a WebAssembly Environment +### Load into a WebAssembly environment Copy the outputs from the build process (e.g. `simple_embedding_vmvx_sync.js` and `simple_embedding_vmvx_sync.wasm`) into your application and follow diff --git a/docs/developers/get_started/cmake_options_and_variables.md b/docs/website/docs/developers/cmake_options_and_variables.md similarity index 71% rename from docs/developers/get_started/cmake_options_and_variables.md rename to docs/website/docs/developers/cmake_options_and_variables.md index 28e6dc6cba15..9c4b137472bc 100644 --- a/docs/developers/get_started/cmake_options_and_variables.md +++ b/docs/website/docs/developers/cmake_options_and_variables.md @@ -1,116 +1,164 @@ -# CMake Options and Variables +# CMake options and variables -## Frequently-used CMake Variables +## Frequently-used CMake variables -### `CMAKE_BUILD_TYPE`:STRING +### `CMAKE_BUILD_TYPE` + +* type: STRING Sets the build type. Possible values are `Release`, `Debug`, `RelWithDebInfo` and `MinSizeRel`. If unset, build type is set to `Release`. -### `CMAKE__COMPILER`:STRING +### `CMAKE__COMPILER` + +* type: STRING This is the command that will be used as the `` compiler, which are `C` and `CXX` in IREE. These variables are set to compile IREE with `clang` or rather `clang++`. Once set, these variables can not be changed. -## IREE-specific CMake Options and Variables +## IREE-specific CMake options and variables This gives a brief explanation of IREE specific CMake options and variables. -### `IREE_ENABLE_RUNTIME_TRACING`:BOOL +### `IREE_ENABLE_RUNTIME_TRACING` + +* type: BOOL Enables instrumented runtime tracing. Defaults to `OFF`. -### `IREE_ENABLE_COMPILER_TRACING`:BOOL +### `IREE_ENABLE_COMPILER_TRACING` + +* type: BOOL Enables instrumented compiler tracing. This requires that `IREE_ENABLE_RUNTIME_TRACING` also be set. Defaults to `OFF`. -### `IREE_BUILD_COMPILER`:BOOL +### `IREE_BUILD_COMPILER` + +* type: BOOL Builds the IREE compiler. Defaults to `ON`. -### `IREE_BUILD_TESTS`:BOOL +### `IREE_BUILD_TESTS` + +* type: BOOL Builds IREE unit tests. Defaults to `ON`. -### `IREE_BUILD_DOCS`:BOOL +### `IREE_BUILD_DOCS` + +* type: BOOL Builds IREE documentation files. Defaults to `OFF`. -### `IREE_BUILD_SAMPLES`:BOOL +### `IREE_BUILD_SAMPLES` + +* type: BOOL Builds IREE sample projects. Defaults to `ON`. -### `IREE_BUILD_PYTHON_BINDINGS`:BOOL +### `IREE_BUILD_PYTHON_BINDINGS` + +* type: BOOL Builds the IREE python bindings. Defaults to `OFF`. -### `IREE_BUILD_BINDINGS_TFLITE`:BOOL +### `IREE_BUILD_BINDINGS_TFLITE` + +* type: BOOL Builds the IREE TFLite C API compatibility shim. Defaults to `ON`. -### `IREE_BUILD_BINDINGS_TFLITE_JAVA`:BOOL +### `IREE_BUILD_BINDINGS_TFLITE_JAVA` + +* type: BOOL Builds the IREE TFLite Java bindings with the C API compatibility shim. Defaults to `ON`. -### `IREE_BUILD_EXPERIMENTAL_REMOTING`:BOOL +### `IREE_BUILD_EXPERIMENTAL_REMOTING` + +* type: BOOL Builds experimental remoting component. Defaults to `OFF`. -### `IREE_HAL_DRIVER_DEFAULTS`:BOOL +### `IREE_HAL_DRIVER_DEFAULTS` + +* type: BOOL Default setting for each `IREE_HAL_DRIVER_*` option. -### `IREE_HAL_DRIVER_*`:BOOL +### `IREE_HAL_DRIVER_*` + +* type: BOOL Individual options enabling the build for each runtime HAL driver. -### `IREE_TARGET_BACKEND_DEFAULTS`:BOOL +### `IREE_TARGET_BACKEND_DEFAULTS` + +* type: BOOL Default setting for each `IREE_TARGET_BACKEND_*` option. -### `IREE_TARGET_BACKEND_*`:BOOL +### `IREE_TARGET_BACKEND_*` + +* type: BOOL Individual options enabling the build for each compiler target backend. -### `IREE_INPUT_*`:BOOL +### `IREE_INPUT_*` + +* type: BOOL Individual options enabling each set of input dialects. -### `IREE_OUTPUT_FORMAT_C`:BOOL +### `IREE_OUTPUT_FORMAT_C` + +* type: BOOL Enables the vm-c compiler output format, using MLIR EmitC. Defaults to `ON`. -### `IREE_DEV_MODE`:BOOL +### `IREE_DEV_MODE` + +* type: BOOL Configure settings to optimize for IREE development (as opposed to CI or release). Defaults to `OFF`. For example, this will downgrade some compiler diagnostics from errors to warnings. -### `IREE_ENABLE_LLD`:BOOL +### `IREE_ENABLE_LLD` + +* type: BOOL Use lld when linking. Defaults to `OFF`. This option is equivalent to `-DIREE_USE_LINKER=lld`. The option `IREE_ENABLE_LLD` and `IREE_USE_LINKER` can not be set at the same time. -### `IREE_ENABLE_ASAN`:BOOL +### `IREE_ENABLE_ASAN` + +* type: BOOL Enable [address sanitizer](https://clang.llvm.org/docs/AddressSanitizer.html) if the current build type is Debug and the compiler supports it. -### `IREE_ENABLE_MSAN`:BOOL +### `IREE_ENABLE_MSAN` + +* type: BOOL Enable [memory sanitizer](https://clang.llvm.org/docs/MemorySanitizer.html) if the current build type is Debug and the compiler supports it. -### `IREE_ENABLE_TSAN`:BOOL +### `IREE_ENABLE_TSAN` + +* type: BOOL Enable [thread sanitizer](https://clang.llvm.org/docs/ThreadSanitizer.html) if the current build type is Debug and the compiler supports it. -### `IREE_ENABLE_UBSAN`:BOOL +### `IREE_ENABLE_UBSAN` + +* type: BOOL Enable [undefiend behavior sanitizer](https://clang.llvm.org/docs/UndefinedBehaviorSanitizer.html) if the current build type is Debug and the compiler supports it. diff --git a/docs/developers/get_started/vulkan_environment_setup.md b/docs/website/docs/developers/vulkan_environment_setup.md similarity index 99% rename from docs/developers/get_started/vulkan_environment_setup.md rename to docs/website/docs/developers/vulkan_environment_setup.md index e06fc33731f0..8d8963444677 100644 --- a/docs/developers/get_started/vulkan_environment_setup.md +++ b/docs/website/docs/developers/vulkan_environment_setup.md @@ -1,4 +1,4 @@ -# Generic Vulkan Development Environment Setup and Troubleshooting +# Vulkan development environment setup and troubleshooting [Vulkan](https://www.khronos.org/vulkan/) is a new generation graphics and compute API that provides high-efficiency, cross-platform access to modern GPUs diff --git a/docs/website/mkdocs.yml b/docs/website/mkdocs.yml index 6249c23d97c1..2b5b206fd1ca 100644 --- a/docs/website/mkdocs.yml +++ b/docs/website/mkdocs.yml @@ -167,6 +167,10 @@ nav: - Extensions: "reference/extensions.md" - "Developers": - "developers/index.md" + - "developers/building_with_bazel.md" + - "developers/building_with_emscripten.md" + - "developers/cmake_options_and_variables.md" + - Vulkan environment setup: "developers/vulkan_environment_setup.md" - "Community": - "community/index.md" - "Blog": From a802979d7851b2d5a3c9772a4889a77672df9e36 Mon Sep 17 00:00:00 2001 From: Scott Todd Date: Wed, 1 Nov 2023 12:04:27 -0700 Subject: [PATCH 03/16] Pull more files across onto the website. --- build_tools/scripts/integrate/README.md | 2 +- .../bazel.md} | 0 .../cmake_options_and_variables.md | 0 .../developers/building/cmake_with_ccache.md} | 2 +- .../emscripten.md} | 0 .../developers/debugging/android_with_lldb.md} | 4 ++-- .../debugging/compile_time_regressions.md | 16 +++++++--------- .../docs/developers/vulkan_environment_setup.md | 4 ++-- docs/website/mkdocs.yml | 14 ++++++++++---- 9 files changed, 23 insertions(+), 19 deletions(-) rename docs/website/docs/developers/{building_with_bazel.md => building/bazel.md} (100%) rename docs/website/docs/developers/{ => building}/cmake_options_and_variables.md (100%) rename docs/{developers/developing_iree/ccache.md => website/docs/developers/building/cmake_with_ccache.md} (98%) rename docs/website/docs/developers/{building_with_emscripten.md => building/emscripten.md} (100%) rename docs/{developers/debugging/lldb_on_android.md => website/docs/developers/debugging/android_with_lldb.md} (98%) rename docs/{ => website/docs}/developers/debugging/compile_time_regressions.md (95%) diff --git a/build_tools/scripts/integrate/README.md b/build_tools/scripts/integrate/README.md index f42de54d03c1..a761d9ffc926 100644 --- a/build_tools/scripts/integrate/README.md +++ b/build_tools/scripts/integrate/README.md @@ -291,7 +291,7 @@ cmake -G Ninja \ ``` To repro failures in CI `bazel_linux_x86-swiftshader_core`, we can follow the -[developer doc](https://iree.dev/developers/bazel) to build IREE using bazel. +[developer doc](https://iree.dev/developers/building/bazel) to build IREE using bazel. E.g., ```bash diff --git a/docs/website/docs/developers/building_with_bazel.md b/docs/website/docs/developers/building/bazel.md similarity index 100% rename from docs/website/docs/developers/building_with_bazel.md rename to docs/website/docs/developers/building/bazel.md diff --git a/docs/website/docs/developers/cmake_options_and_variables.md b/docs/website/docs/developers/building/cmake_options_and_variables.md similarity index 100% rename from docs/website/docs/developers/cmake_options_and_variables.md rename to docs/website/docs/developers/building/cmake_options_and_variables.md diff --git a/docs/developers/developing_iree/ccache.md b/docs/website/docs/developers/building/cmake_with_ccache.md similarity index 98% rename from docs/developers/developing_iree/ccache.md rename to docs/website/docs/developers/building/cmake_with_ccache.md index 203699880228..759b98cc6439 100644 --- a/docs/developers/developing_iree/ccache.md +++ b/docs/website/docs/developers/building/cmake_with_ccache.md @@ -1,4 +1,4 @@ -# Using `ccache` to build IREE +# CMake with `ccache` [`ccache`](https://ccache.dev/) is a compilation cache. In principle, just prepending compiler invocations with `ccache` is all one needs to enable it, diff --git a/docs/website/docs/developers/building_with_emscripten.md b/docs/website/docs/developers/building/emscripten.md similarity index 100% rename from docs/website/docs/developers/building_with_emscripten.md rename to docs/website/docs/developers/building/emscripten.md diff --git a/docs/developers/debugging/lldb_on_android.md b/docs/website/docs/developers/debugging/android_with_lldb.md similarity index 98% rename from docs/developers/debugging/lldb_on_android.md rename to docs/website/docs/developers/debugging/android_with_lldb.md index 1fb6e7d8ef74..d8cb6eb08722 100644 --- a/docs/developers/debugging/lldb_on_android.md +++ b/docs/website/docs/developers/debugging/android_with_lldb.md @@ -1,9 +1,9 @@ +# Android debugging with LLDB + This doc shows how to use LLDB to debug native binaries on Android. For a more complete explanation, see the [official LLDB documentation on remote debugging](https://lldb.llvm.org/use/remote.html). -# Debugging with LLDB on Android - ## Prerequisites We assume the following setup: diff --git a/docs/developers/debugging/compile_time_regressions.md b/docs/website/docs/developers/debugging/compile_time_regressions.md similarity index 95% rename from docs/developers/debugging/compile_time_regressions.md rename to docs/website/docs/developers/debugging/compile_time_regressions.md index d15f1dc23ebe..e88a7335cc1f 100644 --- a/docs/developers/debugging/compile_time_regressions.md +++ b/docs/website/docs/developers/debugging/compile_time_regressions.md @@ -1,4 +1,4 @@ -# Debugging Compile Time Regressions +# Debugging compile time regressions So the IREE compiler used to compile a program quickly, but it is now slower. What do you do? @@ -7,26 +7,26 @@ What do you do? Try to answer as many of these questions as you can: -* **When did compilation get slower?** +> **When did compilation get slower?** A specific git commit is ideal, but "sometime in the last week" is a good starting point. You'll ultimately want to find a culprit release or git commit that changed the compiler code. -* **How much slower did compilation get?** +> **How much slower did compilation get?** Be specific - did it jump from 1 minute to 2 minutes, or 1 minute to 1 hour? Identifying the scale of the regression can help set the priority to investigate it. -* **What is the full compile command?** +> **What is the full compile command?** Try to extract the input program and full list of flags passed to the compiler binary so that others can reproduce what you're seeing. Try to distill this as much as possible to using just native tools (no Python or other framework layers). -* **What environment is the compiler running in?** +> **What environment is the compiler running in?** Are you using a `Debug` build, or a release build? What operating system and size machine is running the compiler (e.g. Linux developer machine, or a @@ -44,7 +44,7 @@ Building the compiler from source and using specific commits in IREE, though it typically won't let you step through changes in submodules (e.g. MLIR updates in `third_party/llvm-project/`). -**Tip**: [Configure ccache](../developing_iree/ccache.md) if you'll be +**Tip**: [Configure ccache](../building/cmake_with_ccache.md) if you'll be rebuilding the compiler while bisecting A manual workflow with `git bisect` looks like this: @@ -81,9 +81,7 @@ git bisect bad [] git bisect run run_bisect.sh ``` -Other sample scripts: - -#### Compile executable sources individually with a timeout +#### Sample: compile executable sources individually with a timeout ```bash #!/bin/bash diff --git a/docs/website/docs/developers/vulkan_environment_setup.md b/docs/website/docs/developers/vulkan_environment_setup.md index 8d8963444677..ff76271c6505 100644 --- a/docs/website/docs/developers/vulkan_environment_setup.md +++ b/docs/website/docs/developers/vulkan_environment_setup.md @@ -1,11 +1,11 @@ -# Vulkan development environment setup and troubleshooting +# Vulkan environment setup [Vulkan](https://www.khronos.org/vulkan/) is a new generation graphics and compute API that provides high-efficiency, cross-platform access to modern GPUs used in a wide variety of devices from PCs and consoles to mobile phones and embedded platforms. -This page lists steps and tips for setting up and trouble shooting a Vulkan +This page lists steps and tips for setting up and troubleshooting a Vulkan development environment. The information here is meant to be generic. ## Vulkan architecture diff --git a/docs/website/mkdocs.yml b/docs/website/mkdocs.yml index 2b5b206fd1ca..264481064e8b 100644 --- a/docs/website/mkdocs.yml +++ b/docs/website/mkdocs.yml @@ -167,10 +167,16 @@ nav: - Extensions: "reference/extensions.md" - "Developers": - "developers/index.md" - - "developers/building_with_bazel.md" - - "developers/building_with_emscripten.md" - - "developers/cmake_options_and_variables.md" - - Vulkan environment setup: "developers/vulkan_environment_setup.md" + - "Building": + - "developers/building/bazel.md" + - "developers/building/emscripten.md" + - "developers/building/cmake_options_and_variables.md" + - "developers/building/cmake_with_ccache.md" + - "Debugging": + - "developers/debugging/android_with_lldb.md" + - "developers/debugging/compile_time_regressions.md" + - "Other topics": + - "developers/vulkan_environment_setup.md" - "Community": - "community/index.md" - "Blog": From d393f4eae956aa79284b19c8dc21b686489b9d1f Mon Sep 17 00:00:00 2001 From: Scott Todd Date: Wed, 1 Nov 2023 13:14:36 -0700 Subject: [PATCH 04/16] Add revision date and document edit links to mkdocs config. --- docs/website/mkdocs.yml | 7 +++++++ docs/website/requirements.txt | 1 + 2 files changed, 8 insertions(+) diff --git a/docs/website/mkdocs.yml b/docs/website/mkdocs.yml index 264481064e8b..d2b1815fa08c 100644 --- a/docs/website/mkdocs.yml +++ b/docs/website/mkdocs.yml @@ -2,6 +2,7 @@ site_name: IREE site_url: https://iree.dev/ repo_url: https://github.com/openxla/iree repo_name: openxla/iree +edit_uri: blob/main/docs/website/docs/ theme: name: material @@ -9,6 +10,7 @@ theme: logo_alt: IREE icon: repo: fontawesome/brands/github + edit: material/file-eye-outline font: text: Noto code: Noto Sans Mono @@ -18,6 +20,7 @@ theme: custom_dir: overrides features: + - content.action.edit # Link to view/edit documentation source on GitHub - content.code.annotate # Allow inline annotations - content.code.copy # Enable copy button - content.tabs.link # Link content tabs across site (e.g. Windows/Linux) @@ -197,6 +200,10 @@ plugins: - tags: tags_file: community/tags.md + # https://squidfunk.github.io/mkdocs-material/setup/adding-a-git-repository/#revisioning + - git-revision-date-localized: + enable_creation_date: false + # https://github.com/mkdocs/mkdocs-redirects - redirects: redirect_maps: # old -> new diff --git a/docs/website/requirements.txt b/docs/website/requirements.txt index a1222cad575b..fbedc4709fb9 100644 --- a/docs/website/requirements.txt +++ b/docs/website/requirements.txt @@ -1,2 +1,3 @@ mkdocs-material==9.2.3 mkdocs-redirects==1.2.1 +mkdocs-git-revision-date-localized-plugin==1.2.1 From 838b7c536c7bd6693bb838d8322260cfa9965712 Mon Sep 17 00:00:00 2001 From: Scott Todd Date: Wed, 1 Nov 2023 13:15:12 -0700 Subject: [PATCH 05/16] Fix document formatting. --- .../website/docs/developers/building/bazel.md | 3 +- .../developers/debugging/android_with_lldb.md | 54 +++++++++---------- .../debugging/compile_time_regressions.md | 4 +- 3 files changed, 32 insertions(+), 29 deletions(-) diff --git a/docs/website/docs/developers/building/bazel.md b/docs/website/docs/developers/building/bazel.md index d2e9f8efd0d1..790345f250c8 100644 --- a/docs/website/docs/developers/building/bazel.md +++ b/docs/website/docs/developers/building/bazel.md @@ -6,7 +6,8 @@ This page walks through building IREE from source using the !!! warning Bazel build support is primarily for internal project infrastructure. We - strongly recommend [using CMake](../building-from-source/index.md) instead. + strongly recommend [using CMake](../../building-from-source/index.md) + instead. Our Bazel configuration is also _only_ tested on Linux. Windows and macOS may be unstable. diff --git a/docs/website/docs/developers/debugging/android_with_lldb.md b/docs/website/docs/developers/debugging/android_with_lldb.md index d8cb6eb08722..7774b7cc2874 100644 --- a/docs/website/docs/developers/debugging/android_with_lldb.md +++ b/docs/website/docs/developers/debugging/android_with_lldb.md @@ -10,9 +10,9 @@ We assume the following setup: 1. [Android NDK is installed](https://developer.android.com/ndk/downloads) and the `ANDROID_NDK` environment variable is set to the installation path. -1. Your Android device connected and configured for +2. Your Android device connected and configured for [`adb`](https://developer.android.com/studio/command-line/adb). -1. The Android binary of interest is already compiled and the command to run it +3. The Android binary of interest is already compiled and the command to run it (in `adb shell`) is ` [program args...]`. This does *not* have to be a proper Android app with a manifest, etc. @@ -20,42 +20,42 @@ We assume the following setup: 1. Push the toolchain files, including `lldb-server`, to your device: - ```shell - adb shell "mkdir -p /data/local/tmp/tools" - adb push "$ANDROID_NDK"/toolchains/llvm/prebuilt/linux-x86_64/lib64/clang/14.0.6/lib/linux/aarch64/* /data/local/tmp/tools - ``` + ```shell + adb shell "mkdir -p /data/local/tmp/tools" + adb push "$ANDROID_NDK"/toolchains/llvm/prebuilt/linux-x86_64/lib64/clang/14.0.6/lib/linux/aarch64/* /data/local/tmp/tools + ``` - You may need to adjust the clang toolchain version to match the one in your - NDK. You can find it with - `find "$ANDROID_NDK/toolchains/llvm/prebuilt" -name lldb-server`. + You may need to adjust the clang toolchain version to match the one in your + NDK. You can find it with + `find "$ANDROID_NDK/toolchains/llvm/prebuilt" -name lldb-server`. -1. Set up port forwarding. We are going to use port 5039 but you are free to +2. Set up port forwarding. We are going to use port 5039 but you are free to pick a different one: - ```shell - adb forward tcp:5039 tcp:5039 - ``` + ```shell + adb forward tcp:5039 tcp:5039 + ``` -1. Start an `lldb-server` in a new interactive adb shell: +3. Start an `lldb-server` in a new interactive adb shell: - ```shell + ```shell adb shell /data/local/tmp/tools/lldb-server platform --listen '*:5039' --server - ``` + ``` -1. Launch `lldb`, connect to the server and run the binary: +4. Launch `lldb`, connect to the server and run the binary: - ```shell - lldb -o 'platform select remote-android' \ + ```shell + lldb -o 'platform select remote-android' \ -o 'platform connect connect://:5039' \ -o 'platform shell cd /data/local/tmp' - target create - run [program args...] - ``` + target create + run [program args...] + ``` - You can either use the system `lldb` or a prebuilt under `"$ANDROID_NDK"/toolchains/llvm/prebuilt/linux-x86_64/lib64/clang/14.0.6/lib/linux/`. + You can either use the system `lldb` or a prebuilt under `"$ANDROID_NDK"/toolchains/llvm/prebuilt/linux-x86_64/lib64/clang/14.0.6/lib/linux/`. - Explanation: each `-o` (short for `--one-shot`) tells lldb to execute a - command on startup. You can run those manually in the lldb shell, if you - prefer. Then, we tell lldb which working directory to use, where to find the - executable, and what command line arguments to use. + Explanation: each `-o` (short for `--one-shot`) tells lldb to execute a + command on startup. You can run those manually in the lldb shell, if you + prefer. Then, we tell lldb which working directory to use, where to find the + executable, and what command line arguments to use. diff --git a/docs/website/docs/developers/debugging/compile_time_regressions.md b/docs/website/docs/developers/debugging/compile_time_regressions.md index e88a7335cc1f..809521d9e893 100644 --- a/docs/website/docs/developers/debugging/compile_time_regressions.md +++ b/docs/website/docs/developers/debugging/compile_time_regressions.md @@ -1,4 +1,4 @@ -# Debugging compile time regressions +# Compile time regression debugging So the IREE compiler used to compile a program quickly, but it is now slower. What do you do? @@ -184,6 +184,8 @@ paint a complete picture and requires waiting for compilation to finish. ### Using Tracy + + See our documentation on [profiling with Tracy](../developing_iree/profiling_with_tracy.md). For compile time regressions, pay particular attention to the different compilation phases From 66f1771bc22f8217303e4087e8f029fffe388814 Mon Sep 17 00:00:00 2001 From: Scott Todd Date: Wed, 1 Nov 2023 13:32:10 -0700 Subject: [PATCH 06/16] Revert revision date (want to configure for generated files). --- docs/website/mkdocs.yml | 4 ---- docs/website/requirements.txt | 1 - 2 files changed, 5 deletions(-) diff --git a/docs/website/mkdocs.yml b/docs/website/mkdocs.yml index d2b1815fa08c..651830f5d520 100644 --- a/docs/website/mkdocs.yml +++ b/docs/website/mkdocs.yml @@ -200,10 +200,6 @@ plugins: - tags: tags_file: community/tags.md - # https://squidfunk.github.io/mkdocs-material/setup/adding-a-git-repository/#revisioning - - git-revision-date-localized: - enable_creation_date: false - # https://github.com/mkdocs/mkdocs-redirects - redirects: redirect_maps: # old -> new diff --git a/docs/website/requirements.txt b/docs/website/requirements.txt index fbedc4709fb9..a1222cad575b 100644 --- a/docs/website/requirements.txt +++ b/docs/website/requirements.txt @@ -1,3 +1,2 @@ mkdocs-material==9.2.3 mkdocs-redirects==1.2.1 -mkdocs-git-revision-date-localized-plugin==1.2.1 From f62dd998f1c1f4f7511eef0e5bbfce71e7664458 Mon Sep 17 00:00:00 2001 From: Scott Todd Date: Wed, 1 Nov 2023 15:34:33 -0700 Subject: [PATCH 07/16] Moving over "general development topics" pages. --- build_tools/scripts/run_markdownlint.sh | 2 - ...integration_correctness_issue_breakdown.md | 51 ---- .../developing_iree/contributing.md | 135 --------- .../developing_iree/contributor_tips.md | 117 -------- docs/developers/developing_iree/releasing.md | 81 ------ .../developing_iree/repository_management.md | 21 -- .../docs}/developers/best_practices.md | 2 +- .../debugging/integration_tests.md} | 47 +++- .../docs}/developers/debugging/releases.md | 2 +- .../docs/developers/debugging}/sanitizers.md | 29 +- .../docs/developers/general/contributing.md | 260 ++++++++++++++++++ .../general/contributing_ci-extra.png} | Bin .../general/contributing_ci_enabled_jobs.png} | Bin .../developers/general}/developer_overview.md | 21 +- .../developers/general/release_management.md | 39 +++ .../developers/general/release_promotion.png} | Bin .../developers/general/release_renaming.png} | Bin .../docs/developers/general}/testing_guide.md | 83 +++--- docs/website/docs/developers/index.md | 4 + docs/website/mkdocs.yml | 10 + 20 files changed, 422 insertions(+), 482 deletions(-) delete mode 100644 docs/developers/debugging/integration_correctness_issue_breakdown.md delete mode 100644 docs/developers/developing_iree/contributing.md delete mode 100644 docs/developers/developing_iree/contributor_tips.md delete mode 100644 docs/developers/developing_iree/releasing.md delete mode 100644 docs/developers/developing_iree/repository_management.md rename docs/{ => website/docs}/developers/best_practices.md (99%) rename docs/{developers/debugging/tf_integrations_test_repro.md => website/docs/developers/debugging/integration_tests.md} (51%) rename docs/{ => website/docs}/developers/debugging/releases.md (99%) rename docs/{developers/developing_iree => website/docs/developers/debugging}/sanitizers.md (82%) create mode 100644 docs/website/docs/developers/general/contributing.md rename docs/{developers/assets/ci-extra.png => website/docs/developers/general/contributing_ci-extra.png} (100%) rename docs/{developers/assets/ci_enabled_jobs.png => website/docs/developers/general/contributing_ci_enabled_jobs.png} (100%) rename docs/{developers/developing_iree => website/docs/developers/general}/developer_overview.md (94%) create mode 100644 docs/website/docs/developers/general/release_management.md rename docs/{developers/assets/promote_release.png => website/docs/developers/general/release_promotion.png} (100%) rename docs/{developers/assets/rename_release.png => website/docs/developers/general/release_renaming.png} (100%) rename docs/{developers/developing_iree => website/docs/developers/general}/testing_guide.md (85%) diff --git a/build_tools/scripts/run_markdownlint.sh b/build_tools/scripts/run_markdownlint.sh index be619b5e6138..0fdb25543349 100755 --- a/build_tools/scripts/run_markdownlint.sh +++ b/build_tools/scripts/run_markdownlint.sh @@ -19,9 +19,7 @@ declare -a included_files_patterns=( "./docs/website/**/*.md" # Some developer documentation .md files that we may move to the website. - "./docs/developers/debugging/*.md" "./docs/developers/developing_iree/*.md" - "./docs/developers/get_started/*.md" ) declare -a excluded_files_patterns=( diff --git a/docs/developers/debugging/integration_correctness_issue_breakdown.md b/docs/developers/debugging/integration_correctness_issue_breakdown.md deleted file mode 100644 index 4996e9f54fa9..000000000000 --- a/docs/developers/debugging/integration_correctness_issue_breakdown.md +++ /dev/null @@ -1,51 +0,0 @@ -This doc describes tips about triaging correctness issue. Feel free to reach out -to @hanhanW or ask questions on Discord if you need help or tips on triaging -correctness issue. - -# Decouple the reproduce from integrations - -## TF integration tests - -See [instructions for reproducing failures in TF/TFLite integration tests](https://github.com/hanhanW/iree/blob/main/docs/developers/debugging/tf_integrations_test_repro.md). - -For input data, they are not dumped within the flagfile. You can construct the -function inputs by looking into `log.txt`. There is an [issue](https://github.com/openxla/iree/issues/8658) -for tracking this. - -## iree-samples - -Follow [README](https://github.com/iree-org/iree-samples#readme) to run the model. -The MLIR files will be generated. You'll find the saved file from log. E.g., - -``` shell -[ RUN ] MobilenetV2Int8Test.test_compile_tflite -I0401 17:27:04.084272 140182373025024 test_util.py:119] Setting up for IREE -I0401 17:27:04.085064 140182373025024 binaries.py:218] Invoke IREE Pipeline: - /tmp/iree-samples/iree-samples.venv/lib/python3.9/site-packages/iree/tools/tflite/iree-import-tflite - /tmp/iree-samples/tflitehub/tmp/mobilenet_v2_int8_test.py/model.tflite - --mlir-print-debuginfo - --save-temp-tfl-input=/tmp/iree-samples/tflitehub/tmp/mobilenet_v2_int8_test.py/tflite.mlir - --save-temp-iree-input=/tmp/iree-samples/tflitehub/tmp/mobilenet_v2_int8_test.py/tosa.mlir -``` - -Unfortunately, the artifacts are not dumped in the runs. There is an [issue](https://github.com/openxla/iree/issues/8756) -for tracking this. A workaround can be found in the issue. - -# Narrow down the repro - -The model itself is big. IREE breaks a model into dispatches and launches the -kernels. The inputs and outputs could be diverged starting from one of -launches. To get a smaller reproduce, you can use [-iree-flow-trace-dispatch-tensors](https://github.com/openxla/iree/blob/main/docs/developers/developing_iree/developer_overview.md#iree-flow-trace-dispatch-tensors). -You can compare the logs between builds/backends, and get the idea about which -dispatch results in wrong outputs. The dumped inputs can be reused in a -flagfile. - -Since we get the suspicious dispatch, we are able to create a test case based on -the dispatch function. The dispatch function can be derived after the -`OutlineDispatchRegions` pass. The function signatures have to be modified -manually. You'll have to put `flow.dispatch.tensor.load` variables to function -arguments, and replace `flow.dispatch.tensor.store` with `return` op. - -At this stage, the reproduce is narrowed down to a single dispatch function. - -Note: This only works when dispatch formation logics are identical between runs. diff --git a/docs/developers/developing_iree/contributing.md b/docs/developers/developing_iree/contributing.md deleted file mode 100644 index d5774eb2538e..000000000000 --- a/docs/developers/developing_iree/contributing.md +++ /dev/null @@ -1,135 +0,0 @@ -# Contributing - -This is a more detailed version of the top-level -[CONTRIBUTING.md](/CONTRIBUTING.md) file. We keep it separate to avoid everyone -getting a pop-up when creating a PR after each time it changes. - -## Downstream Integrations - -Due to the process by which we synchronize this GitHub project with our internal -Google source code repository, there are some oddities in our workflows and -processes. We aim to minimize these, and especially to mitigate their impact on -external contributors, but when they come up, they are documented here for -clarity and transparency. If any of these things are particularly troublesome or -painful for your workflow, please reach out to us so we can prioritize a fix. - -Hopefully these quirks actually make usage in other downstream projects easier, -but integrators may need to look past some details (like the Bazel build system, -Android support, etc.) based on their specific requirements. - -## Build Systems - -IREE supports building from source with both Bazel and CMake. CMake is the -preferred build system for open source users and offers the most flexible -configuration options. Bazel is a stricter build system and helps with usage in -the Google internal source repository. Certain dependencies (think large/complex -projects like CUDA, TensorFlow, PyTorch, etc.) may be difficult to support with -one build system or the other, so the project may configure these as optional. - -## CI - -IREE uses [GitHub Actions](https://docs.github.com/en/actions) for CI. The -primary CI is configured in the -[ci.yml workflow file](/.github/workflows/ci.yml). - -### Self-Hosted Runners - -In addition to the default runners GitHub provides, IREE uses -[self-hosted runners](https://docs.github.com/en/actions/hosting-your-own-runners/managing-self-hosted-runners/about-self-hosted-runners) -to run many of its workflow jobs. These enable access to additional compute and -custom configurations such as accelarators. Configuration scripting is checked -in to this repository (see the -[README for that directory](/build_tools/github_actions/runner/README.md)). - -### Custom Managed Runners - -In addition to our self-hosted runners, we use GitHub's -[large managed runners](https://docs.github.com/en/actions/using-github-hosted-runners/about-larger-runners) -for some platforms that are more trouble to configure ourselves (e.g. Mac). - -### CI Behavior Manipulation - -The setup step of the CI determines which CI jobs to run. This is controlled by -the [configure_ci.py](/build_tools/github_actions/configure_ci.py) script. It -will generally run a pre-determined set of jobs on presubmit with some jobs kept -as post-submit only. If changes are only to a certain set of excluded files that -we know don't affect CI (e.g. docs), then it will skip the jobs. You can -customize which jobs run using -[git trailers](https://git-scm.com/docs/git-interpret-trailers) in the PR -description. The available options are - -``` text -ci-skip: jobs,to,skip -ci-extra: extra,jobs,to,run -ci-exactly: exact,set,of,jobs,to,run -skip-ci: free form reason -skip-llvm-integrate-benchmark: free form reason -benchmark-extra: extra,benchmarks,to,run -runner-env: [testing|prod] -``` - -The first three follow the same format and instruct the setup script on which -jobs to include or exclude from its run. They take a comma-separated list of -jobs which must be from the set of top-level job identifiers in ci.yml file or -the special keyword "all" to indicate all jobs. `ci-skip` removes jobs that -would otherwise be included, though it is not an error to list jobs that would -not be included by default. `ci-extra` adds additional jobs that would not have -otherwise been run, though it is not an error to list jobs that would have been -included anyway. It *is* an error to list a job in both of these fields. -`ci-exactly` provides an exact list of jobs that should run. It is mutually -exclusive with both `ci-skip` and `ci-extra`. In all these cases, the setup does -not make any effort to ensure that job dependencies are satisfied. Thus, if you -request skipping the `build_all` job, all the jobs that depend on it will fail, -not be skipped. `skip-ci` is an older option that simply skips all jobs. Its -usage is deprecated and it is mutually exclusive with all of the other `ci-*` -options. Prefer `ci-skip: all`. - -Benchmarks don't run by default on PRs, and must be specifically requested. They -*do* run by default on PRs detected to be an integration of LLVM into IREE, but -this behavior can be disabled with `skip-llvm-integrate-benchmark`. The -`benchmark-extra` option allows specifying additional benchmark presets to run -as part of benchmarking. It accepts a comma-separated list of benchmark presets. -This combines with labels added to the PR (which are a more limited set of -options). See the [benchmark suites documentation](./benchmark_suites.md). - -The `runner-env` option controls which runner environment to target for our -self-hosted runners. We maintain a test environment to allow testing out new -configurations prior to rolling them out. This trailer is for advanced users who -are working on the CI infrastructure itself. - -#### CI configuration recipes - -Copy/paste any of these at the bottom of a PR description to change what the CI -runs. - -* Also run Windows and macOS builds that are normally post-merge only: - - ``` text - ci-extra: build_test_all_windows,build_test_all_macos_arm64,build_test_all_macos_x86_64 - ``` - -* Also run GPU tests on NVIDIA A100 runners (opt-in due to low availability): - - ``` text - ci-extra: test_a100 - ``` - -* Skip all CI builds and tests, e.g. for comment-only changes: - - ``` text - skip-ci: Comment-only change. - ``` - -* Only run Bazel builds, e.g. for changes only affecting Bazel rules: - - ``` text - ci-exactly: build_test_all_bazel - ``` - -For example, this PR opted in to running the `build_test_all_windows` job: - -![ci-extra](/docs/developers/assets/ci-extra.png) - -The enabled jobs can be viewed from the Summary page of an action run: - -![ci_enabled_jobs](/docs/developers/assets/ci_enabled_jobs.png) diff --git a/docs/developers/developing_iree/contributor_tips.md b/docs/developers/developing_iree/contributor_tips.md deleted file mode 100644 index dd7947a992a5..000000000000 --- a/docs/developers/developing_iree/contributor_tips.md +++ /dev/null @@ -1,117 +0,0 @@ -# Contributor Tips - -This is an opinionated guide documenting workflows that some members of the team -have found useful. It is focused on meta-tooling, not on IREE code specifically -(you will find the latter in the -[Developer Overview](./developer_overview.md)). It is certainly possible to use -workflows other than these, but some common tasks, especially for maintainers -will likely be made easier if you use these flows. It assumes a basic knowledge -of `git` and GitHub and suggests some specific ways of using it. - -## Git Structure - -We tend to use the "triangular" or "forking" workflow. Develop primarily on a -clone of the repository on your development machine. Any local branches named -the same as persistent branches from the -[main repository](https://github.com/openxla/iree) (currently `main` and -`stable`) are pristine (though potentially stale) copies. You only fastforward -these to match upstream and otherwise do development on other branches. When -sending PRs, you push to a different branch on your public fork and create the -PR from there. - -### Setup - -1. Create a fork of the main repository. - -2. Create a local git repository with remotes `upstream` (the main repository) - and `origin` (your personal fork). To list your current remotes - `git remote -v`. - - a. If you already cloned from the main repository (e.g. by following the - getting started guide): - - ```shell - # From your existing git repo - $ git remote rename origin upstream - $ git remote add origin https://github.com//iree.git - ``` - - b. If you haven't already cloned: - - ```shell - # From whatever directory under which you want to nest your repo - $ git clone https://github.com//iree.git - $ cd iree - $ git remote add upstream https://github.com/openxla/iree.git - ``` - - This is especially important for maintainers who have write access (so can - push directly to the main repository) and admins who have elevated - privileges (so can push directly to protected branches). These names are - just suggestions, but you might find some scripts where the defaults are for - remotes named like this. For extra safety, you can make it difficult to push - directly to upstream by setting the push url to something invalid: `git - remote set-url --push upstream DISABLE`, which requires re-enabling the push - URL explicitly before pushing. You can wrap this behavior in a custom git - command like - [git-sudo](https://gist.github.com/GMNGeoffrey/42dd9a9792390094a43bdb69659320c0). - -3. Use a script like - [git_update.sh](/build_tools/scripts/git/git_update.sh) - to easily synchronize `main` with `upstream`. Submodules make this is a - little trickier than it should be. You can also turn this into a git command - by adding it to your path as `git-update`. - -#### Git Config - -These are some additional options you could put in your top-level `.gitconfig` -or repository-specific `.git/config` files that are conducive the recommended -workflow - -```ini -[push] - default = current -[alias] - # Delete branches that you pushed and have been deleted upstream, e.g. because - # the PR was merged. - gone = ! "git fetch -p && git for-each-ref --format '%(refname:short) %(upstream:track)' | awk '$2 == \"[gone]\" {print $1}' | xargs -r git branch -D" - # Update from upstream (custom command) and delete obsolete local branches. - sync = ! (git update main && git gone) - # Create a new branch based off of main (requires a clean working directory). - new = "!f(){ \\\ngit checkout main && git switch -c $1; \\\n}; f" - # Display branches in a useful "latest last" format - br = for-each-ref --sort=committerdate refs/heads/ --format='%(HEAD) %(color:yellow)%(refname:short)%(color:reset) - %(color:red)%(objectname:short)%(color:reset) - %(contents:subject) (%(color:green)%(committerdate:relative)%(color:reset))' - # `git git foo` -> `git foo` typo fixer - git = "!f(){ \\\n git \"$@\"; \\\n}; f" - # Get the git root directory - root = rev-parse --show-toplevel - # checkout, but also sync submodules - ch = "!f() { \\\n git checkout \"$@\"; git submodule sync && git submodule update --init; \\\n}; f" - # See the diff for a PR branch vs the main branch - diffmain = diff --merge-base main - # See only the files that differ vs the main branch - whatsout = diffmain --name-only -[checkout] - # If the checkout command - defaultRemote = origin -[pull] - # When pulling, only complete the pull if its a clean fast forward. - ff = only -[remote] - # Push to your fork (origin) by default - pushDefault = origin -[url "ssh://git@github.com/"] - # Pull with https (so no auth required), but push with ssh. - pushInsteadOf = https://github.com/ -``` - -## Useful Tools - -* GitHub CLI (). A CLI for interacting with GitHub. - Most importantly, it allows scripting the creation of pull requests. -* Refined GitHub Chrome and Firefox Extension: - . Nice extension that adds a - bunch of features to the GitHub UI. -* VSCode: . The most commonly used IDE amongst - IREE developers. -* [ccache](./ccache.md) diff --git a/docs/developers/developing_iree/releasing.md b/docs/developers/developing_iree/releasing.md deleted file mode 100644 index 1e4246273e4f..000000000000 --- a/docs/developers/developing_iree/releasing.md +++ /dev/null @@ -1,81 +0,0 @@ -# Releasing - -IREE cuts automated releases via a workflow that is -[triggered daily](https://github.com/openxla/iree/blob/main/.github/workflows/schedule_candidate_release.yml). -The only constraint placed on the commit that is released is that it has passed -all CI checks. These are published on GitHub with the "pre-release" status. For -debugging this process, see the -[Debugging releases playbook](/docs/developers/debugging/releases.md). - -We periodically promote one of these candidates to a "stable" release by -removing the "pre-release" status. This makes it show up as a "latest" release -on GitHub. We also push the Python packages for this release to PyPI. - -The criteria for selecting this candidate is a bit more involved. - -## Coupling to the Google Integrate - -The Google team that manages these stable releases at the moment is also -responsible for integrating the IREE source code into Google's monorepo. For -convenience, we select a candidate pre-release, attempt to integrate it into -Google's monorepo and then promote it to stable if that was successful without -cherry picks that would affect the quality of the release (because they wouldn't -be present in the promoted version). - -This gives some additional validation to the release because it is stress-tested -running in a different environment and we not-infrequently catch some issues. We -do not currently have a way to add cherry-picks to a release, so if cherry-picks -for functional issues are required, we skip promoting the candidate to "stable". - -This coupling introduces some additional constraints to the process that are not -inherent. It would be perfectly fine to start promoting candidates based on some -other process, but since the same people are managing both, we've coupled these -so we don't have to keep track of as many different versions. - -As the project matures, we will likely remove this coupling. In particular we -will likely start integrating into Google's monorepo at a faster cadence than -the stable releases, so a 1:1 mapping there will not make sense. - -The PyPI password is also currently stored in Google's internal secret -management system, so for others to manage the deployment, we would need to -store it elsewhere with appropriate ACLs. - -At the point where others want to engage in the release process, we can easily -remove any coupling to any Google processes. - -## Picking a Candidate to Promote - -When selecting a candidate we use the following criteria: - -1. ⪆4 days old so that problems with it may have been spotted. -2. Contains no P0 regressions vs the previous stable release. -3. LLVM submodule commit exists upstream (no cherry picks or patches) and - matches a commit already integrated into Google's monorepo - -The constraint on LLVM version is largely due to our current process for doing -so. We aim to lift this limitation and if the process were decoupled from the -Google integration (see -[Coupling to the Google Integrate](#coupling-to-the-google-integrate)), it would -go away anyway. - -There is currently no specific tracking for P0 regressions (process creation in -progress). When you've identified a potential candidate, email the iree-discuss -list with the proposal and solicit feedback. People may point out known -regressions or request that some feature make the cut. - -## Releasing - -1. (Googlers only) Integrate into Google's monorepo, following - . If OSS-relevant cherry-picks were - required to complete this, STOP: do not promote the candidate. - -2. (Googlers only) Push to PyPI using - [pypi_deploy.sh](/build_tools/python_deploy/pypi_deploy.sh) and the - password stored at . - -3. Open the release on GitHub. Rename the release from “candidate" to “stable", - uncheck the option for “pre-release”, and check the option for "latest". - - ![rename_release](/docs/developers/assets/rename_release.png) - - ![promote_release](/docs/developers/assets/promote_release.png) diff --git a/docs/developers/developing_iree/repository_management.md b/docs/developers/developing_iree/repository_management.md deleted file mode 100644 index 3fb3e3aa5707..000000000000 --- a/docs/developers/developing_iree/repository_management.md +++ /dev/null @@ -1,21 +0,0 @@ -# IREE Repository Management - -Due to the process by which we synchronize this GitHub project with our internal -Google source code repository, there are some oddities in our workflows and -processes. We aim to minimize these, and especially to mitigate their impact on -external contributors, but they are documented here for clarity and -transparency. If any of these things are particularly troublesome or painful for -your workflow, please reach out to us so we can prioritize a fix. - -Hopefully these quirks actually make usage in other downstream projects easier, -but integrators may need to look past some details (like the Bazel build system, -Android support, etc.) based on their specific requirements. - -## Build Systems - -IREE supports building from source with both Bazel and CMake. CMake is the -preferred build system for open source users and offers the most flexible -configuration options. Bazel is a stricter build system and helps with usage in -the Google internal source repository. Certain dependencies (think large/complex -projects like CUDA, TensorFlow, PyTorch, etc.) may be difficult to support with -one build system or the other, so the project may configure these as optional. diff --git a/docs/developers/best_practices.md b/docs/website/docs/developers/best_practices.md similarity index 99% rename from docs/developers/best_practices.md rename to docs/website/docs/developers/best_practices.md index 3df6d81718f2..101143bfa527 100644 --- a/docs/developers/best_practices.md +++ b/docs/website/docs/developers/best_practices.md @@ -1,4 +1,4 @@ -# IREE Best Practices +# IREE best practices This page contains a list of best practices for getting the most out of IREE, spanning model authoring, ahead-of-time compilation, and runtime use. Treat diff --git a/docs/developers/debugging/tf_integrations_test_repro.md b/docs/website/docs/developers/debugging/integration_tests.md similarity index 51% rename from docs/developers/debugging/tf_integrations_test_repro.md rename to docs/website/docs/developers/debugging/integration_tests.md index 3cde538df294..26e719a998a7 100644 --- a/docs/developers/debugging/tf_integrations_test_repro.md +++ b/docs/website/docs/developers/debugging/integration_tests.md @@ -1,4 +1,49 @@ -# Debugging failures in TF/TFLite integration tests +# Integration test debugging + +This document includes tips for triaging integration test correctness issues. +Feel free to reach out to @hanhanW or ask questions on Discord for more help. + +## General tips + +### Narrow down reproducers + +* Models themselves can be large, and IREE breaks models into dispatches/kernels +and then launches those individually. Program outputs could diverge starting +from any individual launch. To get a smaller reproducer, you can use +[--iree-flow-trace-dispatch-tensors](../general/developer_overview.md#-iree-flow-trace-dispatch-tensors). +* You can compare the logs between builds/backends to get an idea about which +dispatch results in wrong outputs. The dumped inputs can be reused in a +flagfile. + +Once a suspicious dispatch is identified, we can create a test case based on +the dispatch function. The dispatch function can be derived after the +`OutlineDispatchRegions` pass. The function signatures have to be modified +manually. You'll have to put `flow.dispatch.tensor.load` variables to function +arguments, and replace `flow.dispatch.tensor.store` with `return` op. + +Note: This only works when dispatch formation logics are identical between runs. + +## iree-samples repository tests + +Follow [README](https://github.com/iree-org/iree-samples#readme) to run the model. +The MLIR files will be generated. You'll find the saved file from log. E.g., + +``` shell +[ RUN ] MobilenetV2Int8Test.test_compile_tflite +I0401 17:27:04.084272 140182373025024 test_util.py:119] Setting up for IREE +I0401 17:27:04.085064 140182373025024 binaries.py:218] Invoke IREE Pipeline: + /tmp/iree-samples/iree-samples.venv/lib/python3.9/site-packages/iree/tools/tflite/iree-import-tflite + /tmp/iree-samples/tflitehub/tmp/mobilenet_v2_int8_test.py/model.tflite + --mlir-print-debuginfo + --save-temp-tfl-input=/tmp/iree-samples/tflitehub/tmp/mobilenet_v2_int8_test.py/tflite.mlir + --save-temp-iree-input=/tmp/iree-samples/tflitehub/tmp/mobilenet_v2_int8_test.py/tosa.mlir +``` + +Unfortunately, the artifacts are not dumped in the runs. There is an +[issue](https://github.com/openxla/iree/issues/8756) for tracking this. A +workaround can be found in the issue. + +## TensorFlow integration tests These are steps to reproduce/address failures in TF/TFLite integration tests. These instructions are most stable on Linux, though they may work with a few diff --git a/docs/developers/debugging/releases.md b/docs/website/docs/developers/debugging/releases.md similarity index 99% rename from docs/developers/debugging/releases.md rename to docs/website/docs/developers/debugging/releases.md index 5b05decda32d..82554d8113ef 100644 --- a/docs/developers/debugging/releases.md +++ b/docs/website/docs/developers/debugging/releases.md @@ -1,4 +1,4 @@ -# Debugging releases playbook +# Release debugging playbook ## Tools and Locations diff --git a/docs/developers/developing_iree/sanitizers.md b/docs/website/docs/developers/debugging/sanitizers.md similarity index 82% rename from docs/developers/developing_iree/sanitizers.md rename to docs/website/docs/developers/debugging/sanitizers.md index 95e362a3cfe2..1c9f3423af27 100644 --- a/docs/developers/developing_iree/sanitizers.md +++ b/docs/website/docs/developers/debugging/sanitizers.md @@ -1,4 +1,4 @@ -# Using Address/Memory/Thread Sanitizers +# Sanitizers (ASan/MSan/TSan) [AddressSanitizer](https://clang.llvm.org/docs/AddressSanitizer.html), [MemorySanitizer](https://clang.llvm.org/docs/MemorySanitizer.html) and @@ -14,15 +14,15 @@ They all incur large overhead, so only enable them while debugging. Tool | Detects | Helps debug what? | Slowdown | Memory overhead | Android support ------ | ------- | ----------------- | -------- | --------------- | --------------- -ASan | Out-of-bounds accesses,
Use-after-free,
Use-after-return,
Memory leaks (*), ... | Crashes,
non-deterministic results,
memory leaks (*) | 2x | 3x | Yes +ASan | Out-of-bounds accesses, use-after-free, use-after-return, memory leaks | Crashes, non-deterministic results, memory leaks | 2x | 3x | Yes MSan | Uninitialized memory reads | Non-deterministic results | 3x | ? | Yes TSan | Data races | Many bugs in multi-thread code | 5x-15x | 5x-10x | [No](https://github.com/android/ndk/issues/1171) -Notes: +!!! note -* (*) See [this - documentation](https://clang.llvm.org/docs/AddressSanitizer.html#memory-leak-detection) - on leak detection. It is only enabled by default on some platforms. + See + [this documentation](https://clang.llvm.org/docs/AddressSanitizer.html#memory-leak-detection) + on leak detection. It is only enabled by default on some platforms. ## Support status and how to enable each sanitizer @@ -71,7 +71,7 @@ That ensures that all tests succeed: no test is expected to fail with TSan. If you know what you're doing (i.e. if you are not building targets that internally involve a LLVM/CPU `iree_bytecode_module`), feel free to locally comment out the CMake error and only set `IREE_ENABLE_TSAN`. Also see a -[past attempt](() to relax that CMake +[past attempt](https://github.com/openxla/iree/pull/8966) to relax that CMake validation. ### MSan (MemorySanitizer) @@ -125,9 +125,12 @@ ANDROID_NDK=~/android-ndk-r21d ./build_tools/scripts/android_symbolize.sh < /tmp Where `/tmp/asan.txt` is where you've pasted the raw sanitizer report. -**Tip:** this script will happily just echo any line that isn't a stack frame. -That means you can feed it the whole `ASan` report at once, and it will output a -symbolized version of it. DO NOT run it on a single stack at a time! That is -unlike the symbolizer tool that's being added in NDK r22, and one of the reasons -why we prefer to keep our own script. For more details see [this -comment](https://github.com/android/ndk/issues/753#issuecomment-719719789) +!!! tip + + This script will happily just echo any line that isn't a stack frame. + That means you can feed it the whole `ASan` report at once, and it will + output a symbolized version of it. DO NOT run it on a single stack at a + time! That is unlike the symbolizer tool that's being added in NDK r22, and + one of the reasons why we prefer to keep our own script. For more details + see + [this comment](https://github.com/android/ndk/issues/753#issuecomment-719719789). diff --git a/docs/website/docs/developers/general/contributing.md b/docs/website/docs/developers/general/contributing.md new file mode 100644 index 000000000000..cea66c124b34 --- /dev/null +++ b/docs/website/docs/developers/general/contributing.md @@ -0,0 +1,260 @@ +# Contributing + +This is a more detailed version of the top-level +[CONTRIBUTING.md](https://github.com/openxla/iree/blob/main/CONTRIBUTING.md) +file. We keep it separate to avoid everyone getting a pop-up when creating a PR +after each time it changes. + + + +## Build systems + +IREE supports building from source with both Bazel and CMake. + +* CMake is the preferred build system and offers the most flexible + configuration options +* Bazel is a stricter build system and helps with usage in Google's downstream + source repository +* Certain dependencies (think large/complex projects like CUDA, TensorFlow, + PyTorch, etc.) may be difficult to support with one build system or the + other, so the project may configure these as optional + +## Continuous integration (CI) + +IREE uses [GitHub Actions](https://docs.github.com/en/actions) for CI. The +primary CI is configured in the +[ci.yml workflow file](https://github.com/openxla/iree/blob/main/.github/workflows/ci.yml). + +### Self-hosted runners + +In addition to the default runners GitHub provides, IREE uses +[self-hosted runners](https://docs.github.com/en/actions/hosting-your-own-runners/managing-self-hosted-runners/about-self-hosted-runners) +to run many of its workflow jobs. These enable access to additional compute and +custom configurations such as accelerators. Configuration scripting is checked +in to this repository (see the +[README for that directory](https://github.com/openxla/iree/blob/main/build_tools/github_actions/runner/README.md)). + +### Custom managed runners + +In addition to our self-hosted runners, we use GitHub's +[large managed runners](https://docs.github.com/en/actions/using-github-hosted-runners/about-larger-runners) +for some platforms that are more trouble to configure ourselves (e.g. Mac). + +### CI behavior manipulation + +The setup step of the CI determines which CI jobs to run. This is controlled by +the [configure_ci.py](https://github.com/openxla/iree/blob/main/build_tools/github_actions/configure_ci.py) script. It +will generally run a pre-determined set of jobs on presubmit with some jobs kept +as post-submit only. If changes are only to a certain set of excluded files that +we know don't affect CI (e.g. docs), then it will skip the jobs. You can +customize which jobs run using +[git trailers](https://git-scm.com/docs/git-interpret-trailers) in the PR +description. The available options are + +``` text +ci-skip: jobs,to,skip +ci-extra: extra,jobs,to,run +ci-exactly: exact,set,of,jobs,to,run +skip-ci: free form reason +skip-llvm-integrate-benchmark: free form reason +benchmark-extra: extra,benchmarks,to,run +runner-env: [testing|prod] +``` + +The first three follow the same format and instruct the setup script on which +jobs to include or exclude from its run. They take a comma-separated list of +jobs which must be from the set of top-level job identifiers in ci.yml file or +the special keyword "all" to indicate all jobs. `ci-skip` removes jobs that +would otherwise be included, though it is not an error to list jobs that would +not be included by default. `ci-extra` adds additional jobs that would not have +otherwise been run, though it is not an error to list jobs that would have been +included anyway. It *is* an error to list a job in both of these fields. +`ci-exactly` provides an exact list of jobs that should run. It is mutually +exclusive with both `ci-skip` and `ci-extra`. In all these cases, the setup does +not make any effort to ensure that job dependencies are satisfied. Thus, if you +request skipping the `build_all` job, all the jobs that depend on it will fail, +not be skipped. `skip-ci` is an older option that simply skips all jobs. Its +usage is deprecated and it is mutually exclusive with all of the other `ci-*` +options. Prefer `ci-skip: all`. + +Benchmarks don't run by default on PRs, and must be specifically requested. They +*do* run by default on PRs detected to be an integration of LLVM into IREE, but +this behavior can be disabled with `skip-llvm-integrate-benchmark`. The +`benchmark-extra` option allows specifying additional benchmark presets to run +as part of benchmarking. It accepts a comma-separated list of benchmark presets. +This combines with labels added to the PR (which are a more limited set of +options). See the [benchmark suites documentation](./benchmark_suites.md). + +The `runner-env` option controls which runner environment to target for our +self-hosted runners. We maintain a test environment to allow testing out new +configurations prior to rolling them out. This trailer is for advanced users who +are working on the CI infrastructure itself. + +#### CI configuration recipes + +Copy/paste any of these at the bottom of a PR description to change what the CI +runs. + +* Also run Windows and macOS builds that are normally post-merge only: + + ``` text + ci-extra: build_test_all_windows,build_test_all_macos_arm64,build_test_all_macos_x86_64 + ``` + +* Also run GPU tests on NVIDIA A100 runners (opt-in due to low availability): + + ``` text + ci-extra: test_a100 + ``` + +* Skip all CI builds and tests, e.g. for comment-only changes: + + ``` text + skip-ci: Comment-only change. + ``` + +* Only run Bazel builds, e.g. for changes only affecting Bazel rules: + + ``` text + ci-exactly: build_test_all_bazel + ``` + +For example, this PR opted in to running the `build_test_all_windows` job: + +![ci-extra](./contributing_ci-extra.png) + +The enabled jobs can be viewed from the Summary page of an action run: + +![ci_enabled_jobs](./contributing_ci_enabled_jobs.png) + +## Contributor tips + +These are opinionated tips documenting workflows that some members of the team +have found useful. They are focused on meta-tooling, not on IREE code +specifically (you will find the latter in the +[Developer Overview](./developer_overview.md)). + +!!! note + + It is certainly possible to use workflows other than these. Some common + tasks, especially for maintainers, will likely be made easier by using + these flows. + +We assume a basic knowledge +of `git` and GitHub and suggests some specific ways of using it. + +### Useful tools + +* GitHub CLI (). A CLI for interacting with GitHub. + Most importantly, it allows scripting the creation of pull requests. +* Refined GitHub Chrome and Firefox Extension: + . Nice extension that adds a + bunch of features to the GitHub UI. +* VSCode: . The most commonly used IDE amongst + IREE developers. +* [Ccache](https://ccache.dev/), a fast C/C++ compiler cache. See our + [CMake with `ccache`](../building/cmake_with_ccache.md) page + +### Git structure + +We tend to use the "triangular" or "forking" workflow. Develop primarily on a +clone of the repository on your development machine. Any local branches named +the same as persistent branches from the +[main repository](https://github.com/openxla/iree) are pristine (though +potentially stale) copies. You only fastforward these to match upstream and +otherwise do development on other branches. When sending PRs, you push to a +different branch on your public fork and create the PR from there. + + + +#### Setup + +1. Create a fork of the main repository. + +2. Create a local git repository with remotes `upstream` (the main repository) + and `origin` (your personal fork). To list your current remotes + `git remote -v`. + + a. If you already cloned from the main repository (e.g. by following the + getting started guide): + + ```shell + # From your existing git repo + $ git remote rename origin upstream + $ git remote add origin https://github.com//iree.git + ``` + + b. If you haven't already cloned: + + ```shell + # From whatever directory under which you want to nest your repo + $ git clone https://github.com//iree.git + $ cd iree + $ git remote add upstream https://github.com/openxla/iree.git + ``` + + This is especially important for maintainers who have write access (so can + push directly to the main repository) and admins who have elevated + privileges (so can push directly to protected branches). These names are + just suggestions, but you might find some scripts where the defaults are for + remotes named like this. For extra safety, you can make it difficult to push + directly to upstream by setting the push url to something invalid: `git + remote set-url --push upstream DISABLE`, which requires re-enabling the push + URL explicitly before pushing. You can wrap this behavior in a custom git + command like + [git-sudo](https://gist.github.com/GMNGeoffrey/42dd9a9792390094a43bdb69659320c0). + +3. Use a script like + [git_update.sh](https://github.com/openxla/iree/blob/main/build_tools/scripts/git/git_update.sh) + to easily synchronize `main` with `upstream`. Submodules make this is a + little trickier than it should be. You can also turn this into a git command + by adding it to your path as `git-update`. + +#### Git config + +These are some additional options you could put in your top-level `.gitconfig` +or repository-specific `.git/config` files that are conducive the recommended +workflow + +```ini +[push] + default = current +[alias] + # Delete branches that you pushed and have been deleted upstream, e.g. because + # the PR was merged. + gone = ! "git fetch -p && git for-each-ref --format '%(refname:short) %(upstream:track)' | awk '$2 == \"[gone]\" {print $1}' | xargs -r git branch -D" + # Update from upstream (custom command) and delete obsolete local branches. + sync = ! (git update main && git gone) + # Create a new branch based off of main (requires a clean working directory). + new = "!f(){ \\\ngit checkout main && git switch -c $1; \\\n}; f" + # Display branches in a useful "latest last" format + br = for-each-ref --sort=committerdate refs/heads/ --format='%(HEAD) %(color:yellow)%(refname:short)%(color:reset) - %(color:red)%(objectname:short)%(color:reset) - %(contents:subject) (%(color:green)%(committerdate:relative)%(color:reset))' + # `git git foo` -> `git foo` typo fixer + git = "!f(){ \\\n git \"$@\"; \\\n}; f" + # Get the git root directory + root = rev-parse --show-toplevel + # checkout, but also sync submodules + ch = "!f() { \\\n git checkout \"$@\"; git submodule sync && git submodule update --init; \\\n}; f" + # See the diff for a PR branch vs the main branch + diffmain = diff --merge-base main + # See only the files that differ vs the main branch + whatsout = diffmain --name-only +[checkout] + # If the checkout command + defaultRemote = origin +[pull] + # When pulling, only complete the pull if its a clean fast forward. + ff = only +[remote] + # Push to your fork (origin) by default + pushDefault = origin +[url "ssh://git@github.com/"] + # Pull with https (so no auth required), but push with ssh. + pushInsteadOf = https://github.com/ +``` diff --git a/docs/developers/assets/ci-extra.png b/docs/website/docs/developers/general/contributing_ci-extra.png similarity index 100% rename from docs/developers/assets/ci-extra.png rename to docs/website/docs/developers/general/contributing_ci-extra.png diff --git a/docs/developers/assets/ci_enabled_jobs.png b/docs/website/docs/developers/general/contributing_ci_enabled_jobs.png similarity index 100% rename from docs/developers/assets/ci_enabled_jobs.png rename to docs/website/docs/developers/general/contributing_ci_enabled_jobs.png diff --git a/docs/developers/developing_iree/developer_overview.md b/docs/website/docs/developers/general/developer_overview.md similarity index 94% rename from docs/developers/developing_iree/developer_overview.md rename to docs/website/docs/developers/general/developer_overview.md index fa0f18970677..1dd173c68283 100644 --- a/docs/developers/developing_iree/developer_overview.md +++ b/docs/website/docs/developers/general/developer_overview.md @@ -1,12 +1,9 @@ -# Developer Overview +# Developer overview This guide provides an overview of IREE's project structure and main tools for developers. -**Note: project layout is evolving at the moment, see -** - -## Project Code Layout +## Project code layout * [/compiler/](https://github.com/openxla/iree/blob/main/compiler/): MLIR dialects, LLVM compiler passes, module translation code, etc. @@ -25,7 +22,7 @@ developers. * [/samples/](https://github.com/openxla/iree/blob/main/samples/): Also see the separate repository -## IREE Compiler Code Layout +## IREE compiler code layout * [API/](https://github.com/openxla/iree/tree/main/compiler/src/iree/compiler/API): Public C API @@ -36,7 +33,7 @@ developers. * [InputConversion/](https://github.com/openxla/iree/tree/main/compiler/src/iree/compiler/InputConversion): Conversions from input dialects and preprocessing -## IREE Runtime Code Layout +## IREE runtime code layout * [base/](https://github.com/openxla/iree/blob/main/runtime/src/iree/base/): Common types and utilities used throughout the runtime @@ -55,7 +52,7 @@ developers. Bytecode **V**irtual **M**achine used to work with IREE modules and invoke IREE functions -## Developer Tools +## Developer tools IREE's core compiler accepts programs in supported input MLIR dialects (e.g. `stablehlo`, `tosa`, `linalg`). Import tools and APIs may be used to convert @@ -189,13 +186,11 @@ For example, to inspect the module translated above: ### Useful generic flags -There are a few useful generic flags when working with IREE tools: - #### Read inputs from a file All the IREE tools support reading input values from a file. This is quite -useful for debugging. Use `-help` for each tool to see what the flag to set. The -inputs are expected to be newline-separated. Each input should be either a +useful for debugging. Use `--help` for each tool to see what the flag to set. +The inputs are expected to be newline-separated. Each input should be either a scalar or a buffer. Scalars should be in the format `type=value` and buffers should be in the format `[shape]xtype=[value]`. For example: @@ -204,7 +199,7 @@ should be in the format `[shape]xtype=[value]`. For example: 1x5x3x1xf32=15,14,13,12,11,10,9,8,7,6,5,4,3,2,1 ``` -#### `iree-flow-trace-dispatch-tensors` +#### `--iree-flow-trace-dispatch-tensors` This flag will enable tracing inputs and outputs for each dispatch function. It is easier to narrow down test cases, since IREE breaks a ML workload into diff --git a/docs/website/docs/developers/general/release_management.md b/docs/website/docs/developers/general/release_management.md new file mode 100644 index 000000000000..5e2e2fa89c4c --- /dev/null +++ b/docs/website/docs/developers/general/release_management.md @@ -0,0 +1,39 @@ +# Release management + +IREE cuts automated releases via a workflow that is +[triggered daily](https://github.com/openxla/iree/blob/main/.github/workflows/schedule_candidate_release.yml). +The only constraint placed on the commit that is released is that it has passed +all CI checks. These are published on GitHub with the "pre-release" status. For +debugging this process, see the +[Release debugging playbook](../debugging/releases.md). + +We periodically promote one of these candidates to a "stable" release by +removing the "pre-release" status. This makes it show up as a "latest" release +on GitHub. We also push the Python packages for this release to PyPI. + +## Picking a candidate to promote + +When selecting a candidate we use the following criteria: + +1. ⪆4 days old so that problems with it may have been spotted +2. Contains no P0 regressions vs the previous stable release +3. LLVM submodule commit ideally exists upstream (no cherry picks or patches) + +When you've identified a potential candidate, email the iree-discuss list with +the proposal and solicit feedback. People may point out known regressions or +request that some feature make the cut. + +## Promoting a candidate to stable + +1. (Authorized users only) Push to PyPI using + [pypi_deploy.sh](https://github.com/openxla/iree/blob/main//build_tools/python_deploy/pypi_deploy.sh) + + + * For Googlers, the password is stored at + +2. Open the release on GitHub. Rename the release from "candidate" to "stable", + uncheck the option for "pre-release", and check the option for "latest". + + ![rename_release](./release_renaming.png) + + ![promote_release](./release_promotion.png) diff --git a/docs/developers/assets/promote_release.png b/docs/website/docs/developers/general/release_promotion.png similarity index 100% rename from docs/developers/assets/promote_release.png rename to docs/website/docs/developers/general/release_promotion.png diff --git a/docs/developers/assets/rename_release.png b/docs/website/docs/developers/general/release_renaming.png similarity index 100% rename from docs/developers/assets/rename_release.png rename to docs/website/docs/developers/general/release_renaming.png diff --git a/docs/developers/developing_iree/testing_guide.md b/docs/website/docs/developers/general/testing_guide.md similarity index 85% rename from docs/developers/developing_iree/testing_guide.md rename to docs/website/docs/developers/general/testing_guide.md index c0028e640ad7..90a81692332a 100644 --- a/docs/developers/developing_iree/testing_guide.md +++ b/docs/website/docs/developers/general/testing_guide.md @@ -1,4 +1,4 @@ -# Testing Guide +# Testing guide Like the IREE project in general, IREE tests are divided into a few different components and use different tooling depending on the needs of that component. @@ -11,27 +11,27 @@ components and use different tooling depending on the needs of that component. | | iree_hal_cts_test_suite | CMake | Host/Device | | Core E2E tests | iree_check_test | Bazel/CMake | Host/Device | | | iree_trace_runner_test | Bazel/CMake | Host/Device | -| | iree_single_backend_generated_trace_runner_test | Bazel/CMake | Host/Device | | | iree_generated_trace_runner_test | Bazel/CMake | Host/Device | | | iree_static_linker_test | CMake | Host/Device | There are also more `*_test_suite` targets that groups test targets with the same configuration together. -## Compiler Tests +## Compiler tests Tests for the IREE compilation pipeline are written as lit tests in the same style as MLIR. -By convention, IREE includes tests for printing and parsing of MLIR ops in -`.../IR/test/{OP_CATEGORY}_ops.mlir` files, tests for folding and -canonicalization in `.../IR/test/{OP_CATEGORY}_folding.mlir` files, and tests -for compiler passes and pipelines in other `.../test/*.mlir` files. +By convention, IREE includes tests for -### Running a Test +* printing and parsing of ops in `.../IR/test/{OP_CATEGORY}_ops.mlir` files +* folding and canonicalization in `.../IR/test/{OP_CATEGORY}_folding.mlir` files +* compiler passes and pipelines in other `.../test/*.mlir` files + +### Running a test For the test - +[`iree/compiler/Dialect/VM/Conversion/MathToVM/test/arithmetic_ops.mlir`](https://github.com/openxla/iree/blob/main/compiler/src/iree/compiler/Dialect/VM/Conversion/MathToVM/test/arithmetic_ops.mlir) With CMake, run this from the build directory: @@ -45,7 +45,7 @@ With Bazel, run this from the repo root: bazel test //compiler/src/iree/compiler/Dialect/VM/Conversion/MathToVM/test:arithmetic_ops.mlir.test ``` -### Writing a Test +### Writing a test For advice on writing MLIR compiler tests, see the [MLIR testing guide](https://mlir.llvm.org/getting_started/TestingGuide/). Tests @@ -56,13 +56,13 @@ dialects and passes and doesn't register some unnecessary core ones. As with most parts of the IREE compiler, these should not have a dependency on the runtime. -### Configuring the Build System +### Configuring the build system In the Bazel BUILD file, create a `iree_lit_test_suite` rule. We usually create a single suite that globs all `.mlir` files in the directory and is called "lit". -```bzl +```python load("//iree/build_tools/bazel:iree_lit_test.bzl", "iree_lit_test_suite") iree_lit_test_suite( @@ -76,7 +76,7 @@ iree_lit_test_suite( ``` There is a corresponding CMake function, calls to which will be generated by our -[Bazel to CMake Converter](https://github.com/openxla/iree/tree/main/build_tools/bazel_to_cmake/bazel_to_cmake.py). +[Bazel to CMake converter](https://github.com/openxla/iree/tree/main/build_tools/bazel_to_cmake/bazel_to_cmake.py). ```cmake iree_lit_test_suite( @@ -92,13 +92,13 @@ iree_lit_test_suite( You can also create a test for a single file with `iree_lit_test`. -## Runtime Tests +## Runtime tests Tests for the runtime C++ code use the -[Google Test](https://github.com/google/googletest) testing framework. They +[GoogleTest](https://github.com/google/googletest) testing framework. They should generally follow the style and best practices of that framework. -### Running a Test +### Running a test For the test [`/runtime/src/iree/base/bitfield_test.cc`](https://github.com/openxla/iree/blob/main/runtime/src/iree/base/bitfield_test.cc): @@ -127,7 +127,7 @@ export CTEST_PARALLEL_LEVEL=$(nproc) To use the Vulkan backend as test driver, you may need to select between a Vulkan implementation from SwiftShader and multiple Vulkan-capable hardware devices. This can be done via environment variables. See the -[generic Vulkan setup](../get_started/vulkan_environment_setup.md#useful-environment-variables) +[generic Vulkan setup](../vulkan_environment_setup.md#useful-environment-variables) page for details regarding these variables. For Bazel, you can persist the configuration in `user.bazelrc` to save typing. @@ -143,10 +143,10 @@ test:vknative --test_env="VK_LAYER_PATH=..." Then you can use `bazel test --config=vkswiftshader` to select SwiftShader as the Vulkan implementation. Similarly for other implementations. -### Writing a Test +### Writing a test -For advice on writing tests in the Googletest framework, see the -[Googletest primer](https://github.com/google/googletest/blob/main/docs/primer.md). +For advice on writing tests in the GoogleTest framework, see the +[GoogleTest primer](https://github.com/google/googletest/blob/main/docs/primer.md). Test files for source file `foo.cc` with build target `foo` should live in the same directory with source file `foo_test.cc` and build target `foo_test`. You should `#include` `iree/testing/gtest.h` instead of any of the gtest or gmock @@ -155,14 +155,14 @@ headers. As with all parts of the IREE runtime, these should not have a dependency on the compiler. -### Configuring the Build System +### Configuring the build system In the Bazel BUILD file, create a `cc_test` target with your test file as the source and any necessary dependencies. Usually, you can link in a standard gtest main function. Use `iree/testing:gtest_main` instead of the `gtest_main` that comes with gtest. -```bzl +```python cc_test( name = "arena_test", srcs = ["arena_test.cc"], @@ -194,13 +194,12 @@ There are other more specific test targets, such as `iree_hal_cts_test_suite`, which are designed to test specific runtime support with template configuration and is not supported by Bazel rules. -## IREE Core End-to-End Tests +## IREE core end-to-end (e2e) tests -Here "End-to-End" means from the input accepted by the IREE core compiler -(dialects like TOSA, MHLO, Linalg, and Arithmetic) to execution using the -IREE runtime components. It does not include tests of the integrations with ML -frameworks (e.g. TensorFlow, TensorFlow Lite) or bindings to other languages -(e.g. Python). +Here "end-to-end" means from the input accepted by the IREE core compiler +(dialects like TOSA, StableHLO, Linalg) to execution using the IREE runtime +components. It does not include tests of the integrations with ML frameworks +(e.g. TensorFlow, PyTorch) or bindings to other languages (e.g. Python). We avoid using the more traditional `lit` tests used elsewhere in the compiler for runtime execution tests. Lit tests require running the compiler tools on @@ -224,7 +223,7 @@ cmake --build . --target iree-test-deps ``` To run e2e model tests in -[generated_e2e_model_tests.cmake](/tests/e2e/stablehlo_models/generated_e2e_model_tests.cmake), +[generated_e2e_model_tests.cmake](https://github.com/openxla/iree/tree/main/tests/e2e/stablehlo_models/generated_e2e_model_tests.cmake), because of their dependencies, `-DIREE_BUILD_E2E_TEST_ARTIFACTS=ON` needs to be set when configuring CMake. Also see [IREE Benchmark Suite Prerequisites](/docs/developers/developing_iree/benchmark_suites.md#prerequisites) @@ -233,7 +232,7 @@ for required packages. ### Running a Test For the test - +[`tests/e2e/xla_ops/floor.mlir`](https://github.com/openxla/iree/tree/main/tests/e2e/xla_ops/floor.mlir) compiled for the VMVX target backend and running on the VMVX driver (here they match exactly, but in principle there's a many-to-many mapping from backends to drivers). @@ -253,9 +252,9 @@ bazel test tests/e2e/xla_ops:check_vmvx_local-task_floor.mlir ### Setting test environments Similarly, you can use environment variables to select Vulkan implementations -for running tests as explained in the [Runtime Tests](#runtime-tests) section. +for running tests as explained in the [Runtime tests](#runtime-tests) section. -### Writing a Test +### Writing a test These tests live in `tests/e2e`. A single test consists of a `.mlir` source file specifying an IREE module where each exported function takes no inputs and @@ -297,7 +296,7 @@ an arbitrary SSA-value you can use `util.optimization_barrier`. Next we use this input constant to exercise the runtime feature under test (in this case, just a single floor operation). Finally, we use a check dialect operation to make an assertion about the output. There are a few different -[assertion operations](https://github.com/openxla/iree/tree/main/compiler/src/iree/compiler/Modules/Check). +[assertion operations](../../reference/mlir-dialects/Check.md). Here we use the `expect_almost_eq_const` op: *almost* because we are comparing floats and want to allow for floating-point imprecision, and *const* because we want to compare it to a constant value. This last part is just syntactic sugar @@ -333,12 +332,12 @@ The "module" name for the test suite comes from the default name for an implicit MLIR module. To give the test suite a more descriptive name, use an explicit named top-level module in this file. -### Configuring the Build System +### Configuring the build system A single `.mlir` source file can be turned into a test target with the `iree_check_test` Bazel macro (and corresponding CMake function). -```bzl +```python load("//build_tools/bazel:iree_check_test.bzl", "iree_check_test") iree_check_test( @@ -357,7 +356,7 @@ Usually we want to create a suite of tests across many backends and drivers. This can be accomplished with additional macros. For a single backend/driver pair: -```bzl +```python load("//build_tools/bazel:iree_check_test.bzl", "iree_check_single_backend_test_suite") iree_check_single_backend_test_suite( @@ -370,12 +369,12 @@ iree_check_single_backend_test_suite( This will generate a separate test target for each file in `srcs` with a name following the convention above as well as a Bazel -[test_suite](https://docs.bazel.build/versions/master/be/general.html#test_suite) +[test_suite](https://bazel.build/reference/be/general#test_suite) called "check_vmvx_local-task" that will run all the generated tests. You can also generate suites across multiple pairs: -```bzl +```python load("//build_tools/bazel:iree_check_test.bzl", "iree_check_test_suite") iree_check_test_suite( @@ -399,11 +398,3 @@ our `CMakeLists.txt` file by There are other test targets that generate tests based on template configuraton and platform detection, such as `iree_static_linker_test`. Those targets are not supported by Bazel rules at this point. - -## Binding Tests - -TODO(laurenzo): Explain binding test setup. - -## Integration Tests - -TODO(silvasean): Explain integration test setup. diff --git a/docs/website/docs/developers/index.md b/docs/website/docs/developers/index.md index aedbb284c9c8..eeff4100c7e9 100644 --- a/docs/website/docs/developers/index.md +++ b/docs/website/docs/developers/index.md @@ -1,3 +1,7 @@ # Developers TODO + +* Disclaimer about freshness +* Overview for each section +* Difference between this category and the other website categories diff --git a/docs/website/mkdocs.yml b/docs/website/mkdocs.yml index 651830f5d520..4905a08851f2 100644 --- a/docs/website/mkdocs.yml +++ b/docs/website/mkdocs.yml @@ -139,6 +139,7 @@ nav: - GPU - CUDA: "guides/deployment-configurations/gpu-cuda.md" - GPU - ROCm: "guides/deployment-configurations/gpu-rocm.md" - GPU - Metal: "guides/deployment-configurations/gpu-metal.md" + # TODO(scotttodd): move tips guide to the top level "Developers" section - "Other topics": - Developer tips and tricks: "guides/developer-tips.md" - "Reference": @@ -170,6 +171,11 @@ nav: - Extensions: "reference/extensions.md" - "Developers": - "developers/index.md" + - "General development topics": + - "developers/general/contributing.md" + - "developers/general/developer_overview.md" + - "developers/general/release_management.md" + - "developers/general/testing_guide.md" - "Building": - "developers/building/bazel.md" - "developers/building/emscripten.md" @@ -178,7 +184,11 @@ nav: - "Debugging": - "developers/debugging/android_with_lldb.md" - "developers/debugging/compile_time_regressions.md" + - "developers/debugging/integration_tests.md" + - "developers/debugging/releases.md" + - "developers/debugging/sanitizers.md" - "Other topics": + - "developers/best_practices.md" - "developers/vulkan_environment_setup.md" - "Community": - "community/index.md" From 3c167de3d65251867ad4c496b166dfaf293d7cd2 Mon Sep 17 00:00:00 2001 From: Scott Todd Date: Thu, 2 Nov 2023 12:02:20 -0700 Subject: [PATCH 08/16] Move more files around. They all need cleanup still. --- CONTRIBUTING.md | 2 +- build_tools/scripts/run_markdownlint.sh | 3 - docs/README.md | 28 -- docs/developers/README.md | 4 - .../developing_iree/e2e_benchmarking.md | 404 ------------------ docs/website/README.md | 37 ++ .../developers/design_docs/codegen_passes.md | 0 .../developers/design_docs/cuda_backend.md | 0 .../developers/design_docs/dynamic_shapes.md | 0 .../developers/design_docs/execution_model.md | 0 .../developers/design_docs/function_abi.md | 0 .../design_docs/hal_driver_features.md | 0 .../developers/design_docs/hlo_to_linalg.png | Bin .../design_docs/linalg_to_spirv.png | Bin .../docs}/developers/design_roadmap.md | 213 +-------- ...st_practices.md => iree_best_practices.md} | 0 .../performance}/benchmark_suites.md | 20 +- .../developers/performance}/benchmarking.md | 0 .../docs/developers/performance}/profiling.md | 6 +- .../performance}/profiling_cpu_events.md | 14 +- .../performance/profiling_gpu_vulkan.md} | 2 +- .../performance}/profiling_with_tracy.md | 0 docs/website/mkdocs.yml | 17 +- 23 files changed, 77 insertions(+), 673 deletions(-) delete mode 100644 docs/README.md delete mode 100644 docs/developers/README.md delete mode 100644 docs/developers/developing_iree/e2e_benchmarking.md rename docs/{ => website/docs}/developers/design_docs/codegen_passes.md (100%) rename docs/{ => website/docs}/developers/design_docs/cuda_backend.md (100%) rename docs/{ => website/docs}/developers/design_docs/dynamic_shapes.md (100%) rename docs/{ => website/docs}/developers/design_docs/execution_model.md (100%) rename docs/{ => website/docs}/developers/design_docs/function_abi.md (100%) rename docs/{ => website/docs}/developers/design_docs/hal_driver_features.md (100%) rename docs/{ => website/docs}/developers/design_docs/hlo_to_linalg.png (100%) mode change 100755 => 100644 rename docs/{ => website/docs}/developers/design_docs/linalg_to_spirv.png (100%) mode change 100755 => 100644 rename docs/{ => website/docs}/developers/design_roadmap.md (77%) rename docs/website/docs/developers/{best_practices.md => iree_best_practices.md} (100%) rename docs/{developers/developing_iree => website/docs/developers/performance}/benchmark_suites.md (96%) rename docs/{developers/developing_iree => website/docs/developers/performance}/benchmarking.md (100%) rename docs/{developers/developing_iree => website/docs/developers/performance}/profiling.md (85%) rename docs/{developers/developing_iree => website/docs/developers/performance}/profiling_cpu_events.md (97%) rename docs/{developers/developing_iree/profiling_vulkan_gpu.md => website/docs/developers/performance/profiling_gpu_vulkan.md} (99%) rename docs/{developers/developing_iree => website/docs/developers/performance}/profiling_with_tracy.md (100%) diff --git a/CONTRIBUTING.md b/CONTRIBUTING.md index 993f8394d20d..bcf8ddf70306 100644 --- a/CONTRIBUTING.md +++ b/CONTRIBUTING.md @@ -68,7 +68,7 @@ write access). ## Details For further details, see our -[detailed contributing guide](/docs/developers/developing_iree/contributing.md), which is +[detailed contributing guide](https://iree.dev/developers/contributing/), which is separate to avoid notifying everyone every time it changes. ## Community Guidelines diff --git a/build_tools/scripts/run_markdownlint.sh b/build_tools/scripts/run_markdownlint.sh index 0fdb25543349..cfdbd92ddece 100755 --- a/build_tools/scripts/run_markdownlint.sh +++ b/build_tools/scripts/run_markdownlint.sh @@ -17,9 +17,6 @@ declare -a included_files_patterns=( # .md files for the website. "./docs/website/**/*.md" - - # Some developer documentation .md files that we may move to the website. - "./docs/developers/developing_iree/*.md" ) declare -a excluded_files_patterns=( diff --git a/docs/README.md b/docs/README.md deleted file mode 100644 index 7174d471af29..000000000000 --- a/docs/README.md +++ /dev/null @@ -1,28 +0,0 @@ -# IREE Documentation - -Documentation exclusively for project developers lives in -[`developers/`](developers/), while the source pages and assets for IREE's -user-focused website live in [`website/`](website/). - -Developer documentation should use GitHub-flavored markdown (see -[this guide](https://guides.github.com/features/mastering-markdown/)), while -the website uses [MkDocs](https://www.mkdocs.org/), with the -[Material for MkDocs](https://squidfunk.github.io/mkdocs-material/) theme. - -A high bar should be set for pages published to the website: - -* Page structure should be consistent across sections -* Content should be kept up to date regularly -* Instructions should be portable across environments (e.g. don't - overspecialize on a specific Linux distribution or a particular version of - Visual Studio on Windows) - -When in doubt, the guide at https://developers.google.com/style offers good -instructions. - -Developer documentation _can_ compromise on each of these points. Pages may -also be promoted to website/ after some refinement. - -For more details on how this is set up, see - -* [IREE Website Overview - July 10, 2023](https://docs.google.com/presentation/d/116TyW_aCsPXmmjRYI2tRqpOwDaGNoV8LDC_j9hsMrDk/edit?usp=sharing) diff --git a/docs/developers/README.md b/docs/developers/README.md deleted file mode 100644 index a2b959960fb8..000000000000 --- a/docs/developers/README.md +++ /dev/null @@ -1,4 +0,0 @@ -# IREE Developer Documentation - -Project documentation, developer guides, and other pages not published on -IREE's user-facing website. diff --git a/docs/developers/developing_iree/e2e_benchmarking.md b/docs/developers/developing_iree/e2e_benchmarking.md deleted file mode 100644 index e8f0eaf222fe..000000000000 --- a/docs/developers/developing_iree/e2e_benchmarking.md +++ /dev/null @@ -1,404 +0,0 @@ -# Benchmark IREE and TFLite - - - -We use our end-to-end TensorFlow integration tests to test compilation and -numerical accuracy, and to generate compilation and benchmarking artifacts. -This allows us to validate that our benchmarks are behaving as we expect them -to, and to run them using valid inputs for each model. - -This guide assumes that you can run the tensorflow integration tests. See -[this doc](https://iree.dev/building-from-source/getting-started/#python-bindings) -for more information. That doc also covers writing new tests, which you'll need -to do if you'd like to benchmark a new TensorFlow model. - -## 1. Run IREE's E2E TensorFlow tests to generate the benchmarking artifacts - -```shell -# Continuing from the "Running Python Tests" section of the doc linked above. -# We need not only Python bindings, but also the TensorFlow compiler frontend. -# Self-contained commands from that doc --- skip that if you have already -# completed a build with Python bindings AND TensorFlow compiler frontend. -$ cd iree-build/ # Make and cd into some build directory -$ cmake ../iree -G Ninja \ - -DCMAKE_C_COMPILER=clang \ - -DCMAKE_CXX_COMPILER=clang++ \ - -DIREE_BUILD_PYTHON_BINDINGS=ON -$ cmake --build . -# Also from the Python get-started doc, set this environment variable: -$ export PYTHONPATH=$(pwd)/bindings/python -``` - -```shell -# --target_backends: All tests allow you to specify one or more backends to generate benchmarking artifacts for. -# --artifacts_dir: The default location for these artifacts is under /tmp/iree/modules -# This is a Python 3 program. On some systems, such as Debian derivatives, -# use 'python3' instead of 'python'. -$ python ../iree/integrations/tensorflow/e2e/matrix_ops_static_test.py \ - --target_backends=iree_vmvx -# View the generated artifacts: -$ tree /tmp/iree/modules/MatrixOpsStaticModule/ -``` - -```shell -# Some tests use additional flags to specify features of the Module to test/benchmark: -# --model: the tf.keras.applications model to test -# --data: the dataset (and corresponding image shapes) to create the model for -$ python ../iree/integrations/tensorflow/e2e/keras/applications/applications_test.py \ - --target_backends=iree_vmvx \ - --model=MobileNetV3Small \ - --data=imagenet -# View the generated artifacts: -$ tree /tmp/iree/modules/MobileNetV3Small/ -``` - -Each test/module has a folder with the following artifacts (filtered to only -include those relevant for benchmarking): - -```shell -# Example for a generic module `ModuleName`: -/tmp/iree/modules/ModuleName - ├── iree_vmvx # Or any other IREE backend. - │ └── compiled.vmfb # A flatbuffer containing IREE's compiled code. - └── tflite - ├── module_method_1.tflite - │ # A method on ModuleName compiled to bytes with TFLite, which can - │ # be used by the TFLite's benchmark_model binary. - ├── module_method_2.tflite - ├── ... - └── traces - ├── traced_function_1 - │ └── graph_path - │ # In general, a trace's name does not have to match the name - │ # of the method(s) on the tf.Module that it calls. This file - │ # points to the correct module_method_*.tflite graph file - │ # for TFLite's benchmark_model to use. - ├── traced_function_2 - └── ... - -# Example for MatrixOpsStaticModule: -/tmp/iree/modules/MatrixOpsStaticModule - ├── iree_llvmcpu - │ └── compiled.vmfb - ├── iree_vmvx - │ └──compiled.vmfb - ├── iree_vulkan - │ └── compiled.vmfb - └── tflite - ├── basic_matmul.tflite - ├── matmul_broadcast_singleton_dimension.tflite - ├── matmul_lhs_batch.tflite - ├── matmul_rhs_batch.tflite - └── traces - ├── basic_matmul - │ └── graph_path - ├── matmul_broadcast_singleton_dimension - │ └── graph_path - ├── matmul_lhs_batch - │ └── graph_path - └── matmul_rhs_batch - └── graph_path -``` - -## 2. Benchmarking IREE on desktop - -See also ./benchmarking.md - -Use iree-benchmark-module to benchmark the generated model. For example, to -benchmark a static left-hand-side batched matmul using `MatrixOpsStaticModule` -on VMVX run: - -```shell -$ tools/iree-benchmark-module \ - --module=/tmp/iree/modules/MatrixOpsStaticModule/iree_vmvx/compiled.vmfb \ - --device=local-task \ - --function=matmul_lhs_batch \ - --input=256x64x32xf32=2 \ - --input=32x16xf32=3 - - -``` - -Note that the arguments to `--input` are shapes plus an arbitrary value -to populate a splat. Some more complicated models might have very different -performance characteristics depending on the input data, so this manual -specification will not work well. - -TODO(#6688): Discuss new yaml trace files. - -## 3. Benchmarking TFLite on desktop - -### 3.1 Build TFLite's `benchmark_model` binary - -```shell -# Enter the TensorFlow Bazel workspace. -$ cd third_party/tensorflow/ - -# Build the benchmark_model binary. -$ bazel build --copt=-mavx2 -c opt \ - //tensorflow/lite/tools/benchmark:benchmark_model - -# By default, TFLite/x86 uses various matrix multiplication libraries. -# It is possible to force it to only use Ruy for all matrix multiplications. -# That is the default on ARM but not on x86. This will overwrite the -# previous binary unless you move it. -# -# Note that Ruy takes care of -mavx2 and other AVX extensions internally, -# so this passing this flag here isn't going to make a difference to -# matrix multiplications. However, the rest of TFLite's kernels outside -# of ruy will still benefit from -mavx2. -$ bazel build --copt=-mavx2 -c opt \ - --define=tflite_with_ruy=true \ - //tensorflow/lite/tools/benchmark:benchmark_model - -# The binary can now be found in the following directory: -$ ls bazel-bin/tensorflow/lite/tools/benchmark/ -``` - -### 3.2 Benchmark the model on TFLite - -We pass TFLite the graph generated from the test above (located at the path from -graph_path). It will generate fake inputs for the model. - -Using `MatrixOpsStaticModule`'s left-hand-side batched matmul again as an -example we can run the benchmark as follows: - -```shell -# Run within `third_party/tensorflow/`. -$ ./bazel-bin/tensorflow/lite/tools/benchmark/benchmark_model \ - --graph=$(cat "/tmp/iree/modules/MatrixOpsStaticModule/tflite/traces/matmul_lhs_batch/graph_path") \ - --warmup_runs=1 \ - --num_threads=1 \ - --num_runs=100 \ - --enable_op_profiling=true -``` - -## 4. Benchmarking IREE on Android - -### 4.1 Prepare the benchmarking tools - -IREE only supports compiling to Android with CMake. Documentation on setting up -your environment to cross-compile to Android can be found -[here](https://iree.dev/building-from-source/android/). - -```shell -# After following the instructions above up to 'Build all targets', the -# iree-benchmark-module binary should be in the following directory: -$ ls build-android/tools/ - -# Copy the benchmarking binary to phone. -$ adb push build-android/tools/iree-benchmark-module /data/local/tmp -``` - -### 4.2 Push the IREE's compilation / benchmarking artifacts to the device - -In this example we'll only copy over the files we need to benchmark a single -module on a single backend, but you can easily copy all of the modules over -as well. - -Using `MatrixOpsStaticModule`'s left-hand-side batched matmul again as an -example: - -```shell -# Make a directory for the module/backend pair we want to benchmark. -$ adb shell mkdir -p /data/local/tmp/MatrixOpsStaticModule/iree_vmvx/ - -# Transfer the files. -$ adb push /tmp/iree/modules/MatrixOpsStaticModule/iree_vmvx/* \ - /data/local/tmp/MatrixOpsStaticModule/iree_vmvx/ -``` - -### 4.3 Benchmark the module - -```shell -$ adb shell /data/local/tmp/iree-benchmark-module \ - --module="/data/local/tmp/MatrixOpsStaticModule/iree_vmvx/compiled.vmfb" \ - --device=local-task \ - --function=matmul_lhs_batch \ - --input=256x64x32xf32=2 \ - --input=32x16xf32=3 -``` - -## 5. Benchmarking TFLite on Android - -### 5.1 Prepare the benchmarking tools - -There are three options for getting TFLite's `benchmark_model` binary for -Android. - -The first two are to build it directly, either in a -[`docker` container](https://www.tensorflow.org/lite/guide/build_android#set_up_build_environment_using_docker) -or -[in your own -environment](https://www.tensorflow.org/lite/guide/build_android#set_up_build_environment_without_docker). -To build TensorFlow tools with Android: - -- Run `./configure` under TensorFlow repo. -- Add the following section to the TensorFlow WORKSPACE file. - -``` starlark -android_ndk_repository( - name="androidndk", - path="/full/path/to/android_ndk", -) -``` - -TODO(hanchung): Place the Android setup to somewhere outside IREE, e.g., -TensorFlow. - -Then you can configure the TFLite `benchmark_model` binary in the following -ways: - -```shell -# Build the benchmark_model binary without any add-ons. -# Note that unlike TFLite/x86, TFLite/ARM uses Ruy by default for all -# matrix multiplications (No need to pass tflite_with_ruy), except for some -# matrix*vector products. Below we show how to force using ruy also for that. -$ bazel build -c opt \ - --config=android_arm64 \ - --cxxopt='--std=c++17' \ - //tensorflow/lite/tools/benchmark:benchmark_model - -# Copy the benchmarking binary to phone and allow execution. -$ adb push bazel-bin/tensorflow/lite/tools/benchmark/benchmark_model \ - /data/local/tmp -$ adb shell chmod +x /data/local/tmp/benchmark_model -``` - -```shell -# Build the benchmark_model binary using ruy even for matrix*vector -# products. This is only worth trying in models that are heavy on matrix*vector -# shapes, typically LSTMs and other RNNs. -$ bazel build -c opt \ - --config=android_arm64 \ - --cxxopt='--std=c++17' \ - --copt=-DTFLITE_WITH_RUY_GEMV \ - //tensorflow/lite/tools/benchmark:benchmark_model - -# Rename the binary for comparison with the standard benchmark_model. -$ mv bazel-bin/tensorflow/lite/tools/benchmark/benchmark_model \ - bazel-bin/tensorflow/lite/tools/benchmark/benchmark_model_plus_ruy_gemv -$ adb push bazel-bin/tensorflow/lite/tools/benchmark/benchmark_model_plus_ruy_gemv \ - /data/local/tmp/ -$ adb shell chmod +x /data/local/tmp/benchmark_model_plus_ruy_gemv -``` - -```shell -# Build the benchmark_model binary with flex. -$ bazel build -c opt \ - --config=android_arm64 \ - --cxxopt='--std=c++17' \ - //tensorflow/lite/tools/benchmark:benchmark_model_plus_flex - -# Copy the benchmarking binary to phone and allow execution. -$ adb push bazel-bin/tensorflow/lite/tools/benchmark/benchmark_model_plus_flex \ - /data/local/tmp -$ adb shell chmod +x /data/local/tmp/benchmark_model_plus_flex -``` - -Alternatively, you can download and install the -[Android Benchmark App](https://www.tensorflow.org/lite/performance/measurement#android_benchmark_app). -If you choose to install the app then you'll have to modify the benchmarking -commands below slightly, as shown in -[this example](https://www.tensorflow.org/lite/performance/measurement#run_benchmark). - -### 5.2 Run the benchmark - -```shell -# Copy the data over to the phone. -$ mkdir -p /data/local/tmp/MatrixOpsStaticModule/tflite -$ adb push /tmp/iree/modules/MatrixOpsStaticModule/tflite/* \ - /data/local/tmp/MatrixOpsStaticModule/tflite/ -``` - -```shell -# Benchmark with TFLite. -$ adb shell taskset f0 /data/local/tmp/benchmark_model \ - --graph=/data/local/tmp/MatrixOpsStaticModule/tflite/matmul_lhs_batch.tflite \ - --warmup_runs=1 \ - --num_threads=1 \ - --num_runs=10 \ -``` - -```shell -# Benchmark with TFLite + RUY GEMV -$ adb shell taskset f0 /data/local/tmp/benchmark_model_plus_ruy_gemv \ - --graph=/data/local/tmp/MatrixOpsStaticModule/tflite/matmul_lhs_batch.tflite \ - --warmup_runs=1 \ - --num_threads=1 \ - --num_runs=10 \ -``` - -```shell -# Benchmark with TFLite + Flex. -$ adb shell taskset f0 /data/local/tmp/benchmark_model_plus_flex \ - --graph=/data/local/tmp/MatrixOpsStaticModule/tflite/matmul_lhs_batch.tflite \ - --warmup_runs=1 \ - --num_threads=1 \ - --num_runs=10 \ -``` - -```shell -# Benchmark with TFLite running on GPU. -$ adb shell taskset f0 /data/local/tmp/benchmark_model \ - --graph=/data/local/tmp/MatrixOpsStaticModule/tflite/matmul_lhs_batch.tflite \ - --warmup_runs=1 \ - --num_threads=1 \ - --num_runs=10 \ - --use_gpu=true -``` - -Running benchmark on GPU won't give op profiling. To detailed profiling -information for GPU you can run the following script: - -```shell -# Op profiling on GPU using OpenCL backend. -$ sh tensorflow/lite/delegates/gpu/cl/testing/run_performance_profiling.sh \ - -m /data/local/tmp/MatrixOpsStaticModule/tflite/matmul_lhs_batch.tflite -``` - -Note: You will have to manually specify the TFLite graph that you want to -benchmark, as the `graph_path` file assumes that the graph has not moved. The -name of the `.tflite` graph that you need to benchmark _may_ be different from -the name of the trace that you want to benchmark, but you can use `cat` on -the `graph_path` file to verify the correct `.tflite` filename if you're unsure. - -Tip:
-nbsp;   Sometimes `benchmark_tool` falls back to use CPU even -when `use_gpu` is set. To get more information, you can turn on traces in the -tool by `adb shell setprop debug.tflite.trace 1`. - -### Profile - -There are 2 profilers built into TFLite's `benchmark_model` program. Both of -them impact latencies, so they should only be used to get a breakdown of the -relative time spent in each operator type, they should not be enabled for the -purpose of measuring a latency. - -The first is `enable_op_profiling`. It's based on timestamps before and after -each op. It's a runtime command-line flag taken by `benchmark_model`. Example: - -``` bash -$ adb shell taskset f0 /data/local/tmp/benchmark_model \ - --graph=/data/local/tmp/MatrixOpsStaticModule/tflite/matmul_lhs_batch.tflite \ - --warmup_runs=1 \ - --num_threads=1 \ - --num_runs=10 \ - --enable_op_profiling=true -``` - -The second is `ruy_profiler`. Despite its name, it's available regardless of -whether `ruy` is used for the matrix multiplications. It's a sampling profiler, -which allows it to provide some more detailed information, particularly on -matrix multiplications. It's a build-time switch: - -``` bash -$ bazel build \ - --define=ruy_profiler=true \ - -c opt \ - --config=android_arm64 \ - //tensorflow/lite/tools/benchmark:benchmark_model -``` - -The binary thus built can be run like above, no command-line flag needed. diff --git a/docs/website/README.md b/docs/website/README.md index 6f777b247274..587beb714934 100644 --- a/docs/website/README.md +++ b/docs/website/README.md @@ -38,3 +38,40 @@ Deploy: ```shell mkdocs gh-deploy --remote-name ``` + +## Website sections and authoring tips + +For more details on how this is set up, see +[IREE Website Overview - July 10, 2023](https://docs.google.com/presentation/d/116TyW_aCsPXmmjRYI2tRqpOwDaGNoV8LDC_j9hsMrDk/edit?usp=sharing) +(though note that the website organization has changed since then). + +For documentation language and style, the guide at +https://developers.google.com/style offers good advice. + +### Building from source + +Instructions on how to build the project from source on supported platforms. + +* Focus on instructions that apply to all users, independent of specific + package managers and development styles +* Set developers up for success with good default options +* Explain how interact with the build system and its outputs + +### Guides + +Workflow-oriented guides showing users how to accomplish tasks + +### Reference + +Unopinionated descriptions of system components + +### Developers + +Less structured pages for project development topics + +* Pages may be "promoted" from this category to another category if they are + generally useful to a wide enough range of developers + +### Community (Blog) + +A place to showcase work across the community diff --git a/docs/developers/design_docs/codegen_passes.md b/docs/website/docs/developers/design_docs/codegen_passes.md similarity index 100% rename from docs/developers/design_docs/codegen_passes.md rename to docs/website/docs/developers/design_docs/codegen_passes.md diff --git a/docs/developers/design_docs/cuda_backend.md b/docs/website/docs/developers/design_docs/cuda_backend.md similarity index 100% rename from docs/developers/design_docs/cuda_backend.md rename to docs/website/docs/developers/design_docs/cuda_backend.md diff --git a/docs/developers/design_docs/dynamic_shapes.md b/docs/website/docs/developers/design_docs/dynamic_shapes.md similarity index 100% rename from docs/developers/design_docs/dynamic_shapes.md rename to docs/website/docs/developers/design_docs/dynamic_shapes.md diff --git a/docs/developers/design_docs/execution_model.md b/docs/website/docs/developers/design_docs/execution_model.md similarity index 100% rename from docs/developers/design_docs/execution_model.md rename to docs/website/docs/developers/design_docs/execution_model.md diff --git a/docs/developers/design_docs/function_abi.md b/docs/website/docs/developers/design_docs/function_abi.md similarity index 100% rename from docs/developers/design_docs/function_abi.md rename to docs/website/docs/developers/design_docs/function_abi.md diff --git a/docs/developers/design_docs/hal_driver_features.md b/docs/website/docs/developers/design_docs/hal_driver_features.md similarity index 100% rename from docs/developers/design_docs/hal_driver_features.md rename to docs/website/docs/developers/design_docs/hal_driver_features.md diff --git a/docs/developers/design_docs/hlo_to_linalg.png b/docs/website/docs/developers/design_docs/hlo_to_linalg.png old mode 100755 new mode 100644 similarity index 100% rename from docs/developers/design_docs/hlo_to_linalg.png rename to docs/website/docs/developers/design_docs/hlo_to_linalg.png diff --git a/docs/developers/design_docs/linalg_to_spirv.png b/docs/website/docs/developers/design_docs/linalg_to_spirv.png old mode 100755 new mode 100644 similarity index 100% rename from docs/developers/design_docs/linalg_to_spirv.png rename to docs/website/docs/developers/design_docs/linalg_to_spirv.png diff --git a/docs/developers/design_roadmap.md b/docs/website/docs/developers/design_roadmap.md similarity index 77% rename from docs/developers/design_roadmap.md rename to docs/website/docs/developers/design_roadmap.md index 7220860f9940..de61593528d3 100644 --- a/docs/developers/design_roadmap.md +++ b/docs/website/docs/developers/design_roadmap.md @@ -1,8 +1,4 @@ -# IREE Design Roadmap - - - - +# Design roadmap A not-so-concise walkthrough of various IREE features that are in the design process and planned for future versions. A lot of the questions around how the @@ -14,167 +10,12 @@ and if there's interest please say hi on the [iree-discuss](https://groups.google.com/forum/#!forum/iree-discuss) mailing list. - - -- [IREE Design Roadmap](#iree-design-roadmap) - - [Input Dialects](#input-dialects) - - [Future MLIR XLA HLO Replacement](#future-mlir-xla-hlo-replacement) - - [`linalg`: High-level Hierarchical Optimization](#linalg-high-level-hierarchical-optimization) - - [XLA HLO: Canonicalizations](#xla-hlo-canonicalizations) - - [XLA HLO: Tensor to Primitive Conversion](#xla-hlo-tensor-to-primitive-conversion) - - [Quantization](#quantization) - - [`flow`: Data- and Execution-Flow Modeling](#flow-data--and-execution-flow-modeling) - - [Avoiding Readbacks with `flow.stream`](#avoiding-readbacks-with-flowstream) - - [Threading `flow.stream` through the CFG](#threading-flowstream-through-the-cfg) - - [Predication of `flow.dispatch`](#predication-of-flowdispatch) - - [Deduping `flow.executable`s](#deduping-flowexecutables) - - [Rematerializing CSE'd Expressions](#rematerializing-csed-expressions) - - [Device Placement](#device-placement) - - [`hal`: Hardware Abstraction Layer and Multi-Architecture Executables](#hal-hardware-abstraction-layer-and-multi-architecture-executables) - - [Allow Targets to Specify `hal.interface`s](#allow-targets-to-specify-halinterfaces) - - [Target-specific Scheduling Specialization](#target-specific-scheduling-specialization) - - [Buffer Usage Tracking](#buffer-usage-tracking) - - [Batched Executable Caching and Precompilation](#batched-executable-caching-and-precompilation) - - [Target-aware Executable Compression](#target-aware-executable-compression) - - [Target-aware Constant Compression](#target-aware-constant-compression) - - [Command Buffer Stateful Deduplication](#command-buffer-stateful-deduplication) - - [Resource Timeline](#resource-timeline) - - [Transient Tensor Ringbuffer](#transient-tensor-ringbuffer) - - [Timeline Semaphores on the Module ABI](#timeline-semaphores-on-the-module-abi) - - [GPU-like CPU Scheduling](#gpu-like-cpu-scheduling) - - [`vm`: Lightweight Virtual Machine](#vm-lightweight-virtual-machine) - - [Coroutines for Batching and Cooperative Scheduling](#coroutines-for-batching-and-cooperative-scheduling) - - [Cellular Batching](#cellular-batching) - - [Lowering to LLVM IR](#lowering-to-llvm-ir) - - [Improved Type Support](#improved-type-support) - - [Indirect Command Buffer/On-Accelerator Execution](#indirect-command-bufferon-accelerator-execution) - - +[TOC] ## Input Dialects - - -### Future MLIR XLA HLO Replacement - - - -IREE's current input dialect is the XLA HLO dialect representing operations on -tensors. This was a pragmatic decision based on having HLO already defined and -proof of existing models being lowered to it from Tensorflow, allowing us to -focus on the IREE-specific portions of work. Unfortunately, HLO is tied to -Tensorflow and has many quirks that would not otherwise have been designed had -that not been the case. There are discussions happening about an upstream MLIR -[Tensor Compute Primitives](https://llvm.discourse.group/t/development-of-high-level-tensor-compute-primitives-dialect-s-and-transformations/388/) -dialect that HLO can be lowered into, allowing IREE (and other backends) to -decouple themselves from XLA and be easier to target from frontends. - -### `linalg`: High-level Hierarchical Optimization - - - -It's required that IREE inputs are all in tensor form (and not in-place memref -updates) in order to perform a large majority of the `flow` transformations. -Recent work in the [Linalg](https://mlir.llvm.org/docs/Dialects/Linalg/) dialect -is adding support for operating on value-semantic tensors, meaning that we can -first apply `mhlo` to `linalg` lowerings and any of the transformations -available in Linalg prior to performing our own `flow` lowerings. The advantage -is that Linalg will have much stronger and principled code motion and nested -loop transformation optimizations than is possible on higher-level ops. As not -all operations can be represented as `linalg` ops IREE will be able to ingest a -mix of `linalg`, `std`, and `mhlo` (or its replacement) ops. - -### XLA HLO: Canonicalizations - - - -Very little effort has been applied to `mhlo` optimizations and there are a -significant number of missing folders, canonicalizers, and simple -transformations. Many of these happen in legacy XLA C++ backends; however we -need them in MLIR so that we can make use of dynamic shapes, mixed dialect -inputs, etc. The `tf2xla` bridge work (converting Tensorflow models into the -corresponding `mhlo` ops) is nearing its initial milestones and afterward we -expect more of these missing pieces to be filled in. - -Examples of the optimizations that will greatly benefit IREE (and any other -backend consuming `mhlo`) include: - -- Eliding unneeded transpose, reshape, and broadcast operations. -- Inserting transpose, reshape, and broadcast operations to allow for more - optimal memory access patterns (such as transposing gather input to allow - for memcpy-like transfers instead of column-wise cache-unfriendly accesses). -- Moving operations above broadcasts such that the smallest amount of work is - performed. - -### XLA HLO: Tensor to Primitive Conversion - - - -HLO only operates on tensor values - even for simple scalars - and this presents -a problem when attempting to determine which code should be specified to run on -accelerators vs. what should run on the host. The canonical example is -`mhlo.while`, which as seen in the example below uses scalar tensors for its -loop iteration counter and comparison. - -```mlir -%start = arith.constant dense<1> : tensor -%bound = arith.constant dense<3> : tensor -%res = "mhlo.while"(%start) ( { -^bb0(%count: tensor): - %1 = "mhlo.compare"(%count, %bound) {comparison_direction = "LT"} : (tensor, tensor) -> tensor - "mhlo.return"(%1) : (tensor) -> () -}, { -^bb0(%count: tensor): - %1 = mhlo.add %count, %count : tensor - "mhlo.return"(%1) : (tensor) -> () -}) : (tensor) -> tensor -``` - -A naïve but correct lowering (what's currently in IREE) would perform the -comparison and increment on the device and insert a host readback to see if the -loop should continue: - -```mlir -func @main() -> tensor attributes {iree.reflection = {f = "I1!R6!B3!t6", fv = "1"}} { - %cst = arith.constant dense<1> : tensor - %cst_0 = arith.constant dense<3> : tensor - %cst_1 = arith.constant dense<1> : vector<3xi32> - br ^bb1(%cst : tensor) -^bb1(%2: tensor): // 2 preds: ^bb0, ^bb2 - %3 = flow.ex.stream.fragment(%arg0 = %cst_1 : vector<3xi32>, %arg1 = %2 : tensor, %arg2 = %cst_0 : tensor) -> tensor { - %8 = flow.dispatch @main_ex_dispatch_0::@main_ex_dispatch_0[%arg0 : vector<3xi32>](%arg1, %arg2) : (tensor, tensor) -> tensor - flow.return %8 : tensor - } - %4 = flow.tensor.load %3 : tensor - cond_br %4, ^bb2(%2 : tensor), ^bb3(%2 : tensor) -^bb2(%5: tensor): // pred: ^bb1 - %6 = flow.ex.stream.fragment(%arg0 = %cst_1 : vector<3xi32>, %arg1 = %5 : tensor) -> tensor { - %8 = flow.dispatch @main_ex_dispatch_1::@main_ex_dispatch_1[%arg0 : vector<3xi32>](%arg1) : (tensor) -> tensor - flow.return %8 : tensor - } - br ^bb1(%6 : tensor) -^bb3(%7: tensor): // pred: ^bb1 - return %7 : tensor -} -``` - -Of note is the `flow.tensor.load` op indicating a host readback. Though this -correctly executes the loop it is extremely inefficient. What's desired is for -the loop iterator and condition to all happen on the host, with the iterator -being passed to the loop body as an argument that can be encoded into a command -buffer in future lowering stages. This eliminates host readback and allows for -much larger `flow.stream` sequences, feeding more into the pipeline for the -accelerator. - -Not all source frontends have this issue (misrepresenting simple host -computation as non-dense tensor operations), and our goal is to add a -transformation that heuristically converts `mhlo` ops acting on small tensors to -`std` ops acting on primitive values (`i32`, `index`, etc). - ### Quantization - - It's assumed that any work related to quantization/compression has happened prior to lowering into IREE dialects. Our plan is to use the proposed [Quantization Transforms](https://llvm.discourse.group/t/rfc-a-proposal-for-implementing-quantization-transformations-in-mlir/655) @@ -191,8 +32,6 @@ inference transforms. ## `flow`: Data- and Execution-Flow Modeling - - The `flow` dialect is designed to allow us to extract as much concurrency as possible from a program and partition IR into the scheduling and execution domains. Today we have the IR structure and transformation flow in place but @@ -204,8 +43,6 @@ the allocation you don't make_ ;) ### Avoiding Readbacks with `flow.stream` - - A majority of the readbacks we have today (manifested as `flow.tensor.load.*` ops) will be removed when we have an [HLO tensor->primitive conversion](#xla-hlo-tensor-to-primitive-conversion). @@ -243,8 +80,6 @@ updated within the same stream to entirely remove the host round-trip. ### Threading `flow.stream` through the CFG - - The current `flow.ex.stream.fragment`, as denoted by the `ex`perimental tag, is a temporary implementation designed to get the concept of streams lowered to the HAL dialect. For streams to be effective at modeling larger concurrency scopes @@ -290,8 +125,6 @@ are no external dependencies. ### Predication of `flow.dispatch` - - While the [`flow.stream` threading through the CFG](#threading-flowstream-through-the-cfg) can remove many of the simpler conditional dispatches there will always be some @@ -324,8 +157,6 @@ overhead is reduced. ### Deduping `flow.executable`s - - While still in the `flow` dialect, the executables are target-agnostic. This makes simple IR tree diffing a potential solution to deduplication. Since most of the dispatches originate from the same source-language library calls in input @@ -337,8 +168,6 @@ done rather trivially. ### Rematerializing CSE'd Expressions - - Common subexpression elimination is performed many times during lowering, however there comes a point where the CSE can introduce false dependencies and additional allocations that are otherwise avoidable. For example if a @@ -381,8 +210,6 @@ extremely non-trivial. ### Device Placement - - While still within the `flow` dialect we have the ability to easily split streams and safely shuffle around operations. Target execution backends can opt into such behavior to ensure that device restrictions such as maximum in-flight @@ -395,8 +222,6 @@ on benchmarks, learned traits via ML, etc. ## `hal`: Hardware Abstraction Layer and Multi-Architecture Executables - - As the IREE HAL is designed almost 1:1 with a compute-only Vulkan API many of the techniques classically used in real-time graphics apply. The benefit we have by modeling our usage of such a low-level API in IR is that the normal work - @@ -406,8 +231,6 @@ the full force of an offline compiler against. ### Allow Targets to Specify `hal.interface`s - - The `hal.interface` op specifies the ABI between the scheduler and the device containing the buffer bindings and additional non-buffer data (parameters, shapes, specialization flags, etc). Today a naïve ordering is used uniformly for @@ -428,8 +251,6 @@ will need additional code. ### Target-specific Scheduling Specialization - - Though the `flow` dialect attempts to fuse as many ops as possible into dispatch regions, it's not always possible for all target backends to schedule a region as a single dispatch. A classic example is algorithms like @@ -459,8 +280,6 @@ present. ### Buffer Usage Tracking - - Many explicit hardware APIs require knowing how buffers are used alongside with where they should be located. For example this additional information determines caching policy on buffer accesses (write-through, write-back, etc), visibility @@ -491,8 +310,6 @@ moved and do it a minimum number of times if it is required. ### Batched Executable Caching and Precompilation - - For targets that may require runtime preprocessing of their executables prior to dispatch, such as SPIR-V or MSL, the IREE HAL provides a caching and batch compilation mechanism based on Vulkan's @@ -525,8 +342,6 @@ valid they will be used to avoid compilation. ### Target-aware Executable Compression - - An advantage of representing executable binaries in IR after translation is that we can apply various post-compilation compression and minification techniques while still know precisely where the executable will be used. This is extremely @@ -550,8 +365,6 @@ relatively plug-in way. ### Target-aware Constant Compression - - It's still an area that needs more research but one goal of the IREE design was to enable efficient target- and context-aware compression of large constants (typically model weights/parameters/embeddings). This may mean reusing existing @@ -566,8 +379,6 @@ do the 4-bit to 32-bit decompression, etc. ### Command Buffer Stateful Deduplication - - The IREE HAL - much like Vulkan it is based on - eschews much of the state that traditional APIs have in favor of (mostly) immutable state objects (pipeline layouts, pipeline states, descriptor sets, etc). There are still a few stateful @@ -582,8 +393,6 @@ move such commands out of loop bodies when possible would be helpful. ### Resource Timeline - - A core concept of the IREE scheduler that allows for overlapping in-flight invocations is that of the resource timeline. This identifies module state that can be in use by multiple invocations and assigns timeline milestones denoting @@ -619,8 +428,6 @@ or emit shape logic without data dependencies. ### Transient Tensor Ringbuffer - - (When properly implemented) almost all buffers required during execution never escape the command buffers they are used in or a single VM invocation. We can trivially identify this from the explicit captures of `flow.stream` and @@ -662,8 +469,6 @@ region). ### Timeline Semaphores on the Module ABI - - Functions calls made across modules (either from C++ into the VM, VM->VM, or VM->C++) should be able to define timeline semaphores used to wait and signal on the call. We can do this by making all exports automatically have the semaphores @@ -683,8 +488,6 @@ interrupt watchdog threads, implicit blocking, and other pitfalls. ### GPU-like CPU Scheduling - - One approach to using multiple cores on a CPU is to perform interior parallelization of operations such as OpenMP or library-call-based custom thread pools (gemmlowp). This works when each individual operation is relatively costly @@ -729,8 +532,6 @@ divisor for workgroup count calculations). ## `vm`: Lightweight Virtual Machine - - The VM is designed as a dynamic linkage ABI, stable bytecode representation, and intermediate lowering IR. Many of the optimizations we can perform on it will benefit all use cases (such as when lowering to LLVM IR) by allowing @@ -739,8 +540,6 @@ to perform on arbitrary LLVM IR. ### Coroutines for Batching and Cooperative Scheduling - - One of the largest features currently missing from the VM is coroutines (aka user-mode fiber scheduling). Coroutines are what will allow us to have multiple in-flight invocations into a module - some of which may be waiting on external @@ -801,8 +600,6 @@ hal.device.yield %semaphore4 >= %sem4_target, %semaphore5 >= %sem5_target #### Cellular Batching - - Though coroutines help throughput there is a way we've found to reduce latency that's been documented as [cellular batching](http://madsys.cs.tsinghua.edu.cn/publications/EUROSYS2018-gao.pdf). @@ -852,8 +649,6 @@ research in the long-term. ### Lowering to LLVM IR - - For scenarios where dynamic module loading is not required and entire modules can be compiled into applications we can lower the VM IR to LLVM IR within MLIR's transformation pipeline. Instead of embedding `vm.call` ops that are @@ -876,8 +671,6 @@ size to precisely the functionality used by the module. ### Improved Type Support - - Currently the VM only supports two types: `i32` and `vm.ref`. This is an intentional limitation such that we can determine what is really needed to express the scheduling we perform, with the idea being that such a limited model @@ -894,8 +687,6 @@ inline calculations that are not worth dispatch overhead/synchronization). ### Indirect Command Buffer/On-Accelerator Execution - - Though IREE will use many different tricks such as [predication](#predication-of-flowdispatch) to build deep pipelines there is still the requirement that the command recording and submission happens on the diff --git a/docs/website/docs/developers/best_practices.md b/docs/website/docs/developers/iree_best_practices.md similarity index 100% rename from docs/website/docs/developers/best_practices.md rename to docs/website/docs/developers/iree_best_practices.md diff --git a/docs/developers/developing_iree/benchmark_suites.md b/docs/website/docs/developers/performance/benchmark_suites.md similarity index 96% rename from docs/developers/developing_iree/benchmark_suites.md rename to docs/website/docs/developers/performance/benchmark_suites.md index ebbf7cf0ef78..5d4bd9446b07 100644 --- a/docs/developers/developing_iree/benchmark_suites.md +++ b/docs/website/docs/developers/performance/benchmark_suites.md @@ -1,4 +1,4 @@ -# IREE Benchmark Suites +# Benchmark suites IREE Benchmarks Suites is a collection of benchmarks for IREE developers to track performance improvements/regressions during development. @@ -12,7 +12,7 @@ trigger the benchmark runs. The results will be compared with Information about the definitions of the benchmark suites can be found in the [IREE Benchmark Suites Configurations](/build_tools/python/benchmark_suites/iree/README.md). -## Running Benchmark Suites Locally +## Running benchmark suites locally ### Prerequisites @@ -22,7 +22,7 @@ Install `iree-import-tf` and `iree-import-tflite` in your Python environment and [TFLite Integration](https://iree.dev/guides/ml-frameworks/tflite/)). -### Choose Benchmark Presets +### Choose benchmark presets IREE Benchmark Suites contain many benchmarks for different devices and model sizes, which can take lots of space and time to build all of them. So benchmarks @@ -56,7 +56,7 @@ export EXECUTION_BENCHMARK_PRESETS="cuda,x86_64" export COMPILATION_BENCHMARK_PRESETS="comp-stats" ``` -### Build Benchmark Suites +### Build benchmark suites Configure IREE with `-DIREE_BUILD_E2E_TEST_ARTIFACTS=ON`: @@ -92,7 +92,7 @@ export E2E_TEST_ARTIFACTS_DIR="${IREE_BUILD_DIR?}/e2e_test_artifacts" > TODO(#13683): Each preset should have its own target to further reduce > unnecessary builds -### Run Benchmarks +### Run benchmarks Export the execution benchmark config: @@ -148,7 +148,7 @@ build_tools/benchmarks/run_benchmarks_on_linux.py \ --mode_regex="4-thread" ``` -### Generate Compilation Statistics (Compilation Benchmarks) +### Generate compilation statistics (compilation benchmarks) Export the compilation benchmark config: @@ -171,7 +171,7 @@ build_tools/benchmarks/collect_compilation_statistics.py \ Note that you need to use [Ninja](https://ninja-build.org/) to build the benchmark suites as the tool collects information from its build log. -### Show Execution / Compilation Benchmark Results +### Show execution / compilation benchmark results If you want to generate a comparison report locally, you can use [diff_local_benchmarks.py](/build_tools/benchmarks/diff_local_benchmarks.py) @@ -193,7 +193,7 @@ build_tools/benchmarks/diff_local_benchmarks.py \ > report.md ``` -### Find Compile and Run Commands to Reproduce Benchmarks +### Find compile and run commands to reproduce benchmarks Each benchmark has its benchmark ID in the benchmark suites, you will see a benchmark ID at: @@ -232,7 +232,7 @@ build_tools/benchmarks/benchmark_helper.py dump-cmds \ --benchmark_id="" ``` -### Get Full List of Benchmarks +### Get full list of benchmarks The commands below output the full list of execution and compilation benchmarks, including the benchmark names and their flags: @@ -245,7 +245,7 @@ build_tools/benchmarks/benchmark_helper.py dump-cmds \ --compilation_benchmark_config="${E2E_TEST_ARTIFACTS_DIR?}/comp_config.json" ``` -## Fetching Benchmark Artifacts from CI +## Fetching benchmark Aartifacts from CI ### 1. Find the corresponding CI workflow run diff --git a/docs/developers/developing_iree/benchmarking.md b/docs/website/docs/developers/performance/benchmarking.md similarity index 100% rename from docs/developers/developing_iree/benchmarking.md rename to docs/website/docs/developers/performance/benchmarking.md diff --git a/docs/developers/developing_iree/profiling.md b/docs/website/docs/developers/performance/profiling.md similarity index 85% rename from docs/developers/developing_iree/profiling.md rename to docs/website/docs/developers/performance/profiling.md index 46cb0f2d9e69..11e1ca475bef 100644 --- a/docs/developers/developing_iree/profiling.md +++ b/docs/website/docs/developers/performance/profiling.md @@ -1,4 +1,4 @@ -# Profiling +# Profiling overview IREE [benchmarking](./benchmarking.md) gives us an accurate and reproducible view of program performance at specific levels of granularity. To analyze system @@ -8,7 +8,7 @@ behavior in more depth, there are various ways to ## Tracy Tracy is a profiler that's been used for a wide range of profiling tasks on -IREE. Refer to [profiling_with_tracy.md](./profiling_with_tracy.md). +IREE. Refer to [Profiling with Tracy](./profiling_with_tracy.md). ## Vulkan GPU Profiling @@ -16,7 +16,7 @@ IREE. Refer to [profiling_with_tracy.md](./profiling_with_tracy.md). interactions and Vulkan API usage details. However, information at a finer granularity, especially inside a particular shader dispatch, is missing. To supplement general purpose tools like Tracy, vendor-specific tools can be used. -Refer to [profiling_vulkan_gpu.md](./profiling_vulkan_gpu.md). +Refer to [Profiling GPUs using Vulkan](./profiling_vulkan_gpu.md). ## CPU cache and other CPU event profiling diff --git a/docs/developers/developing_iree/profiling_cpu_events.md b/docs/website/docs/developers/performance/profiling_cpu_events.md similarity index 97% rename from docs/developers/developing_iree/profiling_cpu_events.md rename to docs/website/docs/developers/performance/profiling_cpu_events.md index e352f6b710e4..f3b7cc21f78d 100644 --- a/docs/developers/developing_iree/profiling_cpu_events.md +++ b/docs/website/docs/developers/performance/profiling_cpu_events.md @@ -1,12 +1,12 @@ -# CPU cache and other CPU event profiling +# Profiling CPUs CPUs are able to [record](https://en.wikipedia.org/wiki/Hardware_performance_counter) certain events that may be relevant when investigating the performance of a program. A -common example of such an event is a ["cache -miss"](https://en.wikipedia.org/wiki/CPU_cache#Cache_miss), when the program -tries to access data in memory that isn't already in some CPU cache, causing -that access to be slower than it could otherwise be. +common example of such an event is a +["cache miss"](https://en.wikipedia.org/wiki/CPU_cache#Cache_miss), when the +program tries to access data in memory that isn't already in some CPU cache, +causing that access to be slower than it could otherwise be. Querying and analyzing this data can be useful, but is hard in two distinct ways: @@ -34,7 +34,7 @@ events) and software events from the kernel (such as page faults and context switches). Anyone may use this system call to implement a profiler, but Linux readily offers one, [`perf`](https://perf.wiki.kernel.org/index.php/Main_Page). -### Preserving Artifacts +### Preserving artifacts By default IREE cleans up any temporary files it creates while running. Tools like perf, however, require those files exist even after the process has exited. @@ -45,7 +45,7 @@ the files. This is only needed for the CPU path when using the system loader. export IREE_PRESERVE_DYLIB_TEMP_FILES=1 ``` -### Desktop Linux +### Desktop linux On desktop Linux we can use [`perf`](https://perf.wiki.kernel.org/index.php/Main_Page). It is provided on diff --git a/docs/developers/developing_iree/profiling_vulkan_gpu.md b/docs/website/docs/developers/performance/profiling_gpu_vulkan.md similarity index 99% rename from docs/developers/developing_iree/profiling_vulkan_gpu.md rename to docs/website/docs/developers/performance/profiling_gpu_vulkan.md index 50fb37cc75c6..f5b2861193b9 100644 --- a/docs/developers/developing_iree/profiling_vulkan_gpu.md +++ b/docs/website/docs/developers/performance/profiling_gpu_vulkan.md @@ -1,4 +1,4 @@ -# Vulkan GPU Profiling +# Profiling GPUs using Vulkan [Tracy](./profiling_with_tracy.md) offers great insights into CPU/GPU interactions and Vulkan API usage diff --git a/docs/developers/developing_iree/profiling_with_tracy.md b/docs/website/docs/developers/performance/profiling_with_tracy.md similarity index 100% rename from docs/developers/developing_iree/profiling_with_tracy.md rename to docs/website/docs/developers/performance/profiling_with_tracy.md diff --git a/docs/website/mkdocs.yml b/docs/website/mkdocs.yml index 4905a08851f2..dec6aab05cbc 100644 --- a/docs/website/mkdocs.yml +++ b/docs/website/mkdocs.yml @@ -187,8 +187,23 @@ nav: - "developers/debugging/integration_tests.md" - "developers/debugging/releases.md" - "developers/debugging/sanitizers.md" + - "Performance": + - "developers/performance/benchmarking.md" + - "developers/performance/benchmark_suites.md" + - "developers/performance/profiling.md" + - "developers/performance/profiling_cpu_events.md" + - "developers/performance/profiling_gpu_vulkan.md" + - "developers/performance/profiling_with_tracy.md" + - "Design docs": + - "developers/design_docs/codegen_passes.md" + - "developers/design_docs/cuda_backend.md" + - "developers/design_docs/dynamic_shapes.md" + - "developers/design_docs/execution_model.md" + - "developers/design_docs/function_abi.md" + - "developers/design_docs/hal_driver_features.md" - "Other topics": - - "developers/best_practices.md" + - "developers/design_roadmap.md" + - "developers/iree_best_practices.md" - "developers/vulkan_environment_setup.md" - "Community": - "community/index.md" From 73791c20b33f7e63403144c781616a909a06e2c5 Mon Sep 17 00:00:00 2001 From: Scott Todd Date: Thu, 2 Nov 2023 13:05:58 -0700 Subject: [PATCH 09/16] Prune old pages, fix relative links. --- .../debugging/compile_time_regressions.md | 2 +- .../developers/design_docs/codegen_passes.md | 659 ------------------ .../developers/design_docs/cuda_backend.md | 84 ++- .../{ => design_docs}/design_roadmap.md | 4 +- .../developers/design_docs/dynamic_shapes.md | 166 ----- .../developers/design_docs/function_abi.md | 8 +- .../design_docs/hal_driver_features.md | 120 ---- .../developers/design_docs/hlo_to_linalg.png | Bin 51823 -> 0 bytes ...model.md => invocation_execution_model.md} | 22 +- .../design_docs/linalg_to_spirv.png | Bin 24451 -> 0 bytes .../docs/developers/general/contributing.md | 3 +- .../docs/developers/general/testing_guide.md | 2 +- .../performance/benchmark_suites.md | 10 +- .../docs/developers/performance/profiling.md | 2 +- docs/website/mkdocs.yml | 7 +- 15 files changed, 100 insertions(+), 989 deletions(-) delete mode 100644 docs/website/docs/developers/design_docs/codegen_passes.md rename docs/website/docs/developers/{ => design_docs}/design_roadmap.md (99%) delete mode 100644 docs/website/docs/developers/design_docs/dynamic_shapes.md delete mode 100644 docs/website/docs/developers/design_docs/hal_driver_features.md delete mode 100644 docs/website/docs/developers/design_docs/hlo_to_linalg.png rename docs/website/docs/developers/design_docs/{execution_model.md => invocation_execution_model.md} (96%) delete mode 100644 docs/website/docs/developers/design_docs/linalg_to_spirv.png diff --git a/docs/website/docs/developers/debugging/compile_time_regressions.md b/docs/website/docs/developers/debugging/compile_time_regressions.md index 809521d9e893..53e961a535f2 100644 --- a/docs/website/docs/developers/debugging/compile_time_regressions.md +++ b/docs/website/docs/developers/debugging/compile_time_regressions.md @@ -187,7 +187,7 @@ paint a complete picture and requires waiting for compilation to finish. See our documentation on -[profiling with Tracy](../developing_iree/profiling_with_tracy.md). For compile +[profiling with Tracy](../performance/profiling_with_tracy.md). For compile time regressions, pay particular attention to the different compilation phases (Flow/Stream/HAL), how many times `TranslateExecutablesPass` runs, and if there are outlier passes that take significantly longer to run than others. diff --git a/docs/website/docs/developers/design_docs/codegen_passes.md b/docs/website/docs/developers/design_docs/codegen_passes.md deleted file mode 100644 index 4f69024bb8ad..000000000000 --- a/docs/website/docs/developers/design_docs/codegen_passes.md +++ /dev/null @@ -1,659 +0,0 @@ -# IREE CPU/GPU Code Generation Pipeline - -This document is intended to provide an overview of the codegen pipeline within -IREE used to generate CPU/GPU code. It intends to give an overview of the main -passes used, the objective of the pass, the current implementation, and what it -is expected to achieve in the long term. - -Note that while the code generation pipeline supports dynamic shapes, this work -is very preliminary. The description of this is not covered here. - -## Input to the codegen pipeline - -The input to the code generation pipeline is the module within the -`hal.executable.variant` operation. Functions within this module that do __not__ -have `Visibility::Private` are the *entry point* functions of the dispatch -region. These are the functions that are *invoked* by the IREE runtime. In -addition, each dispatch region also contains a `hal.interface` operation that -describes the ABI to use for the dispatch region. Two examples of the input to -the code generation pipeline are shown below. In both of these, a single -dispatch function contains a sequence of MHLO operations that the dispatch -region creation has grouped into a single region. Ideally the grouped operations -are fused into a single kernel. - -```mlir -hal.executable.variant "vulkan*" { - module attributes {spv.target_env = ...} { - func @main_ex_dispatch() { - %c0 = constant 0 : index - %0 = hal.interface.load.tensor @legacy_io::@arg0, - offset = %c0 : tensor<32x24xf32> - %1 = hal.interface.load.tensor @legacy_io::@arg1, - offset = %c0 : tensor<24x16xf32> - %2 = "mhlo.dot"(%0, %1) {precision_config = ["DEFAULT", "DEFAULT"]} : - (tensor<32x24xf32>, tensor<24x16xf32>) -> tensor<32x16xf32> - hal.interface.store.tensor %2, @legacy_io::@ret0, - offset = %c0 : tensor<32x16xf32> - return - } - hal.interface private @legacy_io { - 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" - } - } -} -``` - - Snippet 1 : Dispatch region with matrix-matrix multiply -operation. - -```mlir -hal.executable.variant "vulkan*" { - module attributes {spv.target_env = ...} { - func @main_ex_dispatch() { - %c0 = constant 0 : index - %0 = hal.interface.load.tensor @legacy_io::@arg0, - offset = %c0 : tensor<10x15xf32> - %1 = hal.interface.load.tensor @legacy_io::@arg1, - offset = %c0 : tensor<10x15xf32> - %2 = hal.interface.load.tensor @legacy_io::@arg2, - offset = %c0 : tensor<15xf32> - %3 = "mhlo.add"(%0, %1) : - (tensor<10x15xf32>, tensor<10x15xf32>) -> tensor<10x15xf32> - %4 = "mhlo.broadcast"(%2) : (tensor<15xf32>) -> tensor<10x15xf32> - %5 = "mhlo.multiply"(%3, %4) : - (tensor<10x15xf32>, tensor<10x15xf32>) -> tensor<10x15xf32> - hal.interface.store.tensor %5, @legacy_io::@ret0, - offset = %c0 : tensor<10x15xf32> - return - } - hal.interface private @legacy_io { - 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 @arg2, set=0, binding=2, - type="StorageBuffer", access="Read" - hal.interface.binding @ret0, set=0, binding=3, - type="StorageBuffer", access="Write|Discard" - } - } -} -``` - - Snippet 2 : Dispatch region with element-wise -operations. - -__Roadmap Note__: The current implementation might not actually fuse the -operations grouped into a dispatch region into a single kernel. It is possible -to end up with multiple kernels per dispatch region. Over time we plan to -address this by using fusion at different levels (see below). - -The inputs to the dispatch region are materialized within the entry point -function using the `hal.interface.load.tensor` operation, This operation returns -a `tensor` view of the buffer used to store the inputs. Similarly the result of -the dispatch region are *written* out using the `hal.interface.store.tensor` -operation. - -The main constraint that the code generation operates under is that it should -not require additional (temporary) buffers to execute the operations grouped -together within a dispatch region. The rationale behind this constraint is that -buffer allocation/synchronization in IREE happens at the granularity of dispatch -regions, allowing the scheduler to make better decision about where to insert -appropriate synchronizations. - -The IR after all the passes used in the lowering from MHLO to SPIR-V for the -above two examples can be found here ([matrix-matrix multiply op][DotAfterAll], -[elementwise ops][PwAfterAll]). Below is a description of the major passes used. - -## Conversion from MHLO dialect to Linalg on buffers - -The code generation pipeline heavily relies on use of -[Structured Operations][LinalgRationale], specifically the -[Linalg Dialect][LinalgDialect]. Both, the Linalg operations on `tensor`s and on -`memref`s are central to the progressive lowering approach followed here. The -first part of the code generation pipeline is to convert the MHLO operations on -`tensor`s to Linalg operation on `memref`s. This part of the pipeline is common -to both CPU and GPU code generation. - -The steps involved in this conversion is shown below. Each of the arrows -represents a pass in the pipeline: - -![MHLO To Linalg on `memref` conversion](./hlo_to_linalg.png) - -The next sections describe each of these passes in more detail. - -### MHLO to Linalg on tensors - -The first step is to convert MHLO operations to Linalg on tensors. This is done -using the [HLOToLinalgPass][HLOToLinalgPass] from Tensorflow. An example of the -conversion is shown below, where the `mhlo.add`, `mhlo.broadcast` and -`mhlo.multiply` operations are converted to `linalg.generic` operations on -tensors. - -```mlir -#map0 = affine_map<(d0, d1) -> (d0, d1)> -#map1 = affine_map<(d0, d1) -> (d1)> -%3 = linalg.generic - {args_in = 2 : i64, args_out = 1 : i64, - indexing_maps = [#map0, #map0, #map0], - iterator_types = ["parallel", "parallel"]} %0, %1 { - ^bb0(%arg0: f32, %arg1: f32): // no predecessors - %5 = addf %arg0, %arg1 : f32 - linalg.yield %5 : f32 - } : tensor<10x15xf32>, tensor<10x15xf32> -> tensor<10x15xf32> -%4 = linalg.generic - {args_in = 1 : i64, args_out = 1 : i64, - indexing_maps = [#map1, #map0], - iterator_types = ["parallel", "parallel"]} %2 { - ^bb0(%arg0: f32): // no predecessors - linalg.yield %arg0 : f32 - }: tensor<15xf32> -> tensor<10x15xf32> -%5 = linalg.generic - {args_in = 2 : i64, args_out = 1 : i64, - indexing_maps = [#map0, #map0, #map0], - iterator_types = ["parallel", "parallel"]} %3, %4 { - ^bb0(%arg0: f32, %arg1: f32): // no predecessors - %5 = mulf %arg0, %arg1 : f32 - linalg.yield %5 : f32 - }: tensor<10x15xf32>, tensor<10x15xf32> -> tensor<10x15xf32> -``` - - Snippet 3 : MHLO to Linalg conversion for -[element-wise operations](#snippet2) - -At the time of writing the representation of Linalg on `tensor`s does not model -reduction iterator types completely. Specifically, the reduction in Linalg is -modeled using read-modify-write approach, i.e. each iteration of the reduction -loop reads the value stored in the output, adds its contribution, and writes -back to the same location. This means the output has to be *initialized* to the -null element of the reduction operator (i.e. 0 if the reduction is done using -addition). This works for operations on buffers. Since tensors are SSA values -they cannot be updated in-place. As a result, the reduction semantics does not -map as well to `tensor`s. For now it is treated as a convention that when the -Linalg operation is converted to use `memref`s it has to be initialized -appropriately before performing the reduction. Due to this, the conversion from -MHLO op to Linalg op is only done for operations which do not need a *reduction* -iterator type in the converted Linalg op. Consequently, only element-wise -operations, broadcast operations and data movement operations (like copy and -transpose) are converted to Linalg operations at this stage. - -__Roadmap note__: One long term solution for the above is to have operations on -tensors that have *reduction* iterator type to take an additional argument that -contains the initial value of the result tensor. When the operation is converted -to use `memref`s, the buffer for the initial value operand can be reused for the -result. The details involved have not been fully worked out yet. - -### Fusion of Linalg on tensor operations - -The Linalg on `tensor` operations generated at the previous step are fused using -the [LinalgFusionOfTensorOps][LinalgFusionOfTensorOps] from MLIR. Since -`tensor`s are SSA values, fusion at this stage can be done without using alias -analysis or dependence analysis based on reads and writes. Instead the use-def -chains for the `tensor` values can be used to implement producer-consumer -fusion. This stage fuses most elementwise operations, broadcast operations and -data movement operations. An example of the fused op is shown below. - -```mlir -#map0 = affine_map<(d0, d1) -> (d0, d1)> -#map1 = affine_map<(d0, d1) -> (d1)> -%3 = linalg.generic - {args_in = 3 : i64, args_out = 1 : i64, - indexing_maps = [#map0, #map0, #map1, #map0], - iterator_types = ["parallel", "parallel"]} %0, %1, %2 { - ^bb0(%arg0: f32, %arg1: f32, %arg2: f32): // no predecessors - %4 = addf %arg0, %arg1 : f32 - %5 = mulf %4, %arg2 : f32 - linalg.yield %5 : f32 - }: tensor<10x15xf32>, tensor<10x15xf32>, tensor<15xf32> -> tensor<10x15xf32> -``` - - Snippet 4: Fusion of Linalg operation on tensors for -element-wise operations shown in [Snippet 3](#snippet3) - -### Conversion of Linalg on tensors to Linalg on buffers - -Post fusion all the operation on `tensor`s are converted to analogous operations -on `memref`s. In general, this requires a buffer allocation pass. In IREE, -buffer allocation happens at the granularity of dispatch region, and as -mentioned [earlier](#input-to-the-codegen-pipeline), the dispatch region is not -expected to use any additional temporary buffers. So instead of having another -buffer allocation pass within the code generation pipeline, a simpler approach -is used within IREE: - -- For each `hal.interface.store.tensor` an `iree.placeholder` operation is - created. The latter uses the same `hal.interface.binding` as the former, but - returns a `memref` view of the output of the dispatch region instead of a - `tensor` view. This `iree.placeholder` operation is added to start of the - entry point function. - -- A map is constructed that for a given `tensor` records the `memref` value to - use during the conversion. In this map the `tensor` value used in the - `hal.interface.store.tensor` is mapped to the `memref` value returned by the - created `iree.placeholder` operation. - -- The Dialect Conversion framework is used to implement a set of patterns that - convert from operations on `tensor`s to operation on `memref`s, - - - A `hal.interface.load.tensor`, is replaced with an `iree.placeholder` to - get the `memref` view of the input to the dispatch region. - - All Linalg operation on `tensor`s (expected to be just `linalg.generic` - or `linalg.indexed_generic` operations) are converted to the - corresponding operation on `memref`s. Instead of returning a `tensor` - value the converted operation takes an additional `memref` operand as - argument. This `memref` is where the result of the operation is - populated. Current implementation looks for the `memref` to use from the - map constructed previously. If there is no `memref` associated with the - result `tensor` the conversion fails. - - At this stage, any `mhlo` operation not converted to a Linalg operation - are directly converted to a Linalg operation on buffers. This is done - for operations that when converted to Linalg have a *reduction* iterator - type. Some examples of ops converted this way are - - - `mhlo.dot` - - `mhlo.reduce` - - `mhlo.conv` - - `mhlo.reduce_window`. - - Since the specification of the Linalg operations require the output - `memref` to be initialized appropriately, a `linalg.fill` operation is - used to achieve this. - -__Roadmap Note__ : Right now the code-generation pipeline relies on fusion of -operations on tensor level. In the near future, we want to be able to fuse -operations like `linalg.matmul` and `linalg.conv` with consumers/producers that -are element-wise operations using the -[fusion of Linalg operation on `memref`s][LinalgFusionOnBuffers]. - -At this stage of the compilation all operations must have been converted to -Linalg operations on buffers. Shown below are the IR at the end of this stage -for the two examples in Snippets 1 and 2. - -```mlir -func @main_ex_dispatch() { - %0 = iree.placeholder for "interface buffer" - {binding = @legacy_io::@ret0} : memref<32x16xf32> - %c0 = constant 0 : index - %1 = iree.placeholder for "interface buffer" - {binding = @legacy_io::@arg0} : memref<32x24xf32> - %2 = iree.placeholder for "interface buffer" - {binding = @legacy_io::@arg1} : memref<24x16xf32> - %cst = constant 0.000000e+00 : f32 - linalg.matmul(%1, %2, %0) : - memref<32x24xf32>, memref<24x16xf32>, memref<32x16xf32> - return -} -``` - - Snippet 5 : Matrix-matrix multiply after conversion to -Linalg operation on `memref`s. - -```mlir -#map0 = affine_map<(d0, d1) -> (d0, d1)> -#map1 = affine_map<(d0, d1) -> (d1)> -func @main_ex_dispatch() { - %0 = iree.placeholder for "interface buffer" - {binding = @legacy_io::@ret0} : memref<10x15xf32> - %c0 = constant 0 : index - %1 = iree.placeholder for "interface buffer" - {binding = @legacy_io::@arg0} : memref<10x15xf32> - %2 = iree.placeholder for "interface buffer" - {binding = @legacy_io::@arg1} : memref<10x15xf32> - %3 = iree.placeholder for "interface buffer" - {binding = @legacy_io::@arg2} : memref<15xf32> - linalg.generic - {args_in = 3 : i64, args_out = 1 : i64, - indexing_maps = [#map0, #map0, #map1, #map0], - iterator_types = ["parallel", "parallel"]} %1, %2, %3, %0 { - ^bb0(%arg0: f32, %arg1: f32, %arg2: f32, %arg3: f32): // no predecessors - %4 = addf %arg0, %arg1 : f32 - %5 = mulf %4, %arg2 : f32 - linalg.yield %5 : f32 - }: memref<10x15xf32>, memref<10x15xf32>, memref<15xf32>, memref<10x15xf32> - return -} -``` - - Snippet 6 : Elementwise operations after conversion to -Linalg operation on `memref`s - -The rest of the code-generation differs on whether the compilation is for CPU -(using LLVM) or for GPU (using SPIR-V). - -## Conversion from Linalg on buffers to SPIR-V dialect - -The following sections describe the progressive lowering of Linalg operation on -buffers to SPIR-V dialect. Once lowered to the SPIR-V dialect, it can be -serialized into a SPIR-V binary using the -[serialization mechanism provided by the SPIR-V dialect][SpirvSerialization]. -The steps involved in the lowering are described below, with each of the arrows -representing a pass. - -![Linalg on `memref` to SPIR-V conversion](./linalg_to_spirv.png) - -These passes are described below in more detail. - -### Tiling and fusion on buffer operations - -The GPU hardware typically provides multiple-levels of compute hierarchy, namely -*workgroup* level, *subgroup* level and *workitem* level. These map to blocks, -warps and threads, respectively, in CUDA terminology. Tiling is a way to map the -computations to each level of the compute hierarchy. For example 3-D tiling a -`linalg.matmul` operation decomposes the computation into several tiled -matrix-matrix multiplies. -[Tiling transformation in Linalg dialect][LinalgTiling] generates the -outer-loops that iterate over tiled `linalg.matmul` operations. These outer -loops can be mapped to different workgroups, if they are parallel. The tiled -`linalg.matmul` operation can be further tiled to map to subgroups. Finally, the -tiled operation can be lowered to loops with individual iterations mapped to -workitems. The [LinalgTileAndFusePass][LinalgTileAndFuse] uses the Linalg Tiling -patterns ([defined here][LinalgTilingPatterns]) to tile operations like -`linalg.matmul`, `linalg.conv` and `linalg.*_pooling`. The result of tiling the -code in Snippet 5 is shown below. As expected there are 2-parallel loops that -iterate over tiles of the original iteration space (i.e. inter-tile loops) and -can be distributed to workgroups. - -```mlir -func @main_ex_dispatch_0() - attributes { - spv.entry_point_abi = {local_size = dense<[8, 8, 1]> : vector<3xi32>}} { - %cst = constant 0.000000e+00 : f32 - %c32 = constant 32 : index - %c24 = constant 24 : index - %c16 = constant 16 : index - %c0 = constant 0 : index - %c4 = constant 4 : index - %0 = iree.placeholder for "interface buffer" - {binding = @legacy_io::@ret0} : memref<32x16xf32> - %1 = iree.placeholder for "interface buffer" - {binding = @legacy_io::@arg0} : memref<32x24xf32> - %2 = iree.placeholder for "interface buffer" - {binding = @legacy_io::@arg1} : memref<24x16xf32> - linalg.fill(%cst, %0) : f32, memref<32x16xf32> - scf.parallel (%arg0, %arg1) = (%c0, %c0) to (%c32, %c16) step (%c8, %c8) { - scf.for %arg2 = %c0 to %24 step %c4 { - ... - %5 = subview %1[%arg0, %arg2]... - ... - %8 = subview %2[%arg2, %arg1]... - ... - %11 = subview %0[%arg0, %arg1].. - linalg.matmul {__internal_linalg_transform__ = "workgroup"} %5, %8, %11... - } - scf.yield - } - return -} -``` - - Snippet 7 : `linalg.matmul` after tiling. - -#### Tile Size and Workgroup Size - -When operations that are to be tiled exist within the dispatch function (like -`linalg.matmul` or `linalg.conv`), this pass also decides the 1. Tile size to be -used for the tiling. 1. The workgroup size to be used. - -The tile size and workgroup size are closely linked since the code within the -tiled loops are to be collectively executed by the entire workgroup. In other -words, all workitems in the workgroup collaborate to execute the tiled -`linalg.matmul`. - -__Roadmap Note__ : Currently the tile sizes used in this pass are hard-wired. -Not much effort has been put into finding ideal tile size for each operation on -different hardware. The value used is meant to be a baseline to test -functionality, with performance considerations addressed over time. - -#### Markers - -Downstream passes have to handle tiled Linalg operations and untiled Linalg -operation that might exist in the same function in different ways. For example, -while the former are to be executed collectively by workitems within a -workgroup, the latter have to be executed by all workitems across workgroups. -One way to distinguish these two operations is to use the marker mechanism in -Linalg ([LinalgTransformationFilter][LinalgTilingPatterns]). This is a `StrAttr` -whose value can be used to encode the scope of the operation. For example, in -Snippet 7 above, the tiled `linalg.matmul` operation has a marker `workgroup` to -indicate that this operation needs to be executed by a workgroup in a collective -manner. At this time, the code-generation pipeline uses only the `workgroup` -marker. - -__Roadmap Note__ : Markers are meant to be short-lived, ideally set and consumed -within the same pass. In the current pipeline the lifetime spans passes to allow -lowering to different hierarchies. The separate passes that implement the -lowering from Linalg to SPIR-V can be combined into a single pass, relying A -> -B -> C translation mechanism of the Dialect Conversion framework to implement -the progressive lowering. In interest of separation of concerns and for better -debuggability these passes are kept separate at the cost of having lifetimes of -markers span passes. - -#### Promoting subviews to use workgroup local memory and use of synchronizations - -`Workgroup` memory (or `shared memory` in CUDA terminology) can be used to -prefetch the inputs to the tiled operation. For example in the matrix-matrix -multiply case, the same data row (column) of the LHS (RHS) matrix is read by -multiple workitems. Prefetching the data into `Workgroup` memory can reduce the -number of loads to `StorageClass` memory by an order of magnitude. This -transformation can be achieved by using the -[`Linalg Promotion`][LinalgPromotionPatterns] which modifies the `subview`s that -are the operands to the tiled Linalg operation to use a new `memref` object. The -size of this `memref` is computed from the size of the `subview`. This `memref` -object is later lowered to use `Workgroup` memory Storage Class. The snippet -below shows this transformation when applied to `linalg.matmul` (along with -tiling). The newly created `memref` objects are annotated with the memory space -`3` to indicate that they are to be lowered to use `Workgroup` memory. The copy -of data from the original `memref` into the new `memref`, as well as the -necessary synchronization constructs are generated as well. Note the memory -space annotation used here is consistent with what -[address space annotations used in NVVM][NVVMAddressSpace]. - -```mlir -func @matmul_tile() - attributes { - spv.entry_point_abi = {local_size = dense<[8, 8, 1]> : vector<3xi32>}} { - %c32 = constant 32 : index - %c24 = constant 24 : index - %c16 = constant 16 : index - %c4 = constant 4 : index - %c8 = constant 8 : index - %c0 = constant 0 : index - %c1 = constant 1 : index - %0 = iree.placeholder for "interface buffer" - {binding = @legacy_io::@arg0} : memref<32x24xf32> - %1 = iree.placeholder for "interface buffer" - {binding = @legacy_io::@arg1} : memref<24x16xf32> - %2 = iree.placeholder for "interface buffer" - {binding = @legacy_io::@ret0} : memref<32x16xf32> - scf.parallel (%arg0, %arg1) = (%c0, %c0) to (%c32, %c16) step (%c8, %c8) { - scf.for %arg2 = %c0 to %c24 step %c4 { - ... - %5 = subview %0[%arg0, %arg2]... - ... - %8 = subview %1[%arg2, %arg1]... - ... - %11 = subview %2[%arg0, %arg1]... - %12 = alloc(%c8, %c4) : memref - %13 = subview %12[%c0, %c0]... - %14 = alloc(%c4, %c8) : memref - %15 = subview %14[%c0, %c0]... - linalg.copy(%5, %13) {__internal_linalg_transform__ = "workgroup"} - : memref, memref - spv.ControlBarrier "Workgroup", "Workgroup", "AcquireRelease" - linalg.copy(%8, %15) {__internal_linalg_transform__ = "workgroup"} - : memref, memref - spv.ControlBarrier "Workgroup", "Workgroup", "AcquireRelease" - linalg.matmul {__internal_linalg_transform__ = "workgroup"} %13, %15, %11... - spv.ControlBarrier "Workgroup", "Workgroup", "AcquireRelease" - dealloc %12 : memref - dealloc %14 : memref - } - scf.yield - } - return -} -``` - - Snippet 8: `linalg.matmul` after tiling and promotion of -operand subviews to use `Workgroup` memory. - -### Distributing to workgroups and workitems - -After tiling the operations within the dispatch functions are either -`scf.parallel` operations or Linalg operations. - -- The outer `scf.parallel` operations represent parallel loops that are to be - distributed across workgroups. The distribution here assumes that the number - of workgroups along each dimension is equal to the number of iterations of - the `scf.parallel` operation. - -- Linalg operations that are not tiled, and are therefore __not within__ `scf` - operations, are lowered to loops. The resulting outer `scf.parallel` - operations are collapsed to have a single induction variable. This loop is - then distributed across workitems using their `GlobalInvocationId`, (which - is same as `blockIdx * blockDim + threadIdx` in CUDA terminology). - -- Linalg operations that are tiled, and are therefore __within__ `scf` - operations, are lowered to loops and the iterations of the `scf.parallel` - operations are mapped to workitems using their `LocalInvocationId` (which is - same as `threadIdx` in CUDA terminology). Note that these operations are - tagged with the `workgroup` marker which makes it easy to disambiguate from - the case where Linalg operations are outside of `scf` operations. Here too, - the distribution assumes that the workgroup size is greater than or equal to - the number of iterations of the partitioned loop. - -These transformations are applied by the [`ConvertToGPUPass`][ConvertToGPU]. -Below is the result of applying this pass to Snippet 7. The outer `scf.parallel` -loop is distributed across workgroups. The tiled `linalg.matmul` operation is -lowered to loops, and the outer `scf.parallel` operation generated during this -lowering are distributed across workitems within the workgroup. - -```mlir -func @main_ex_dispatch_0_dispatch_1() - attributes { - spv.entry_point_abi = {local_size = dense<[8, 8, 1]> : vector<3xi32>}} { - %c24 = constant 24 : index - %c8 = constant 8 : index - %c4 = constant 4 : index - %c0 = constant 0 : index - %c1 = constant 1 : index - %0 = iree.placeholder for "interface buffer" - {binding = @legacy_io::@ret0} : memref<32x16xf32> - %1 = iree.placeholder for "interface buffer" - {binding = @legacy_io::@arg0} : memref<32x24xf32> - %2 = iree.placeholder for "interface buffer" - {binding = @legacy_io::@arg1} : memref<24x16xf32> - %3 = "gpu.block_id"() {dimension = "x"} : () -> index - %4 = "gpu.block_id"() {dimension = "y"} : () -> index - %5 = muli %4, %c8 : index - %6 = muli %3, %c8 : index - scf.for %arg0 = %c0 to %c24 step %c4 { - ... - %15 = subview %1[%5, %arg0] - ... - %20 = subview %2[%arg0, %6] - %21 = subview %0[%5, %6] - %22 = "gpu.thread_id"() {dimension = "x"} : () -> index - %23 = "gpu.thread_id"() {dimension = "y"} : () -> index - %24 = cmpi "slt", %23, %10 : index - %25 = cmpi "slt", %22, %19 : index - %26 = and %24, %25 : i1 - scf.if %26 { - scf.for %arg1 = %c0 to %14 step %c1 { - %27 = load %15[%23, %arg1] : memref - %28 = load %20[%arg1, %22] : memref - %29 = load %21[%23, %22] : memref - %30 = mulf %21, %22 : f32 - %31 = addf %23, %24 : f32 - store %25, %15[%23, %22] : memref<4x?xf32, #map1> - } - } - } - return -} -``` - - Snippet 9: `linalg.matmul` after distributing parallel -inter-tile loops to workgroups and intra-tile loops to workitems. - -[Snippet 6](#snippet6) shows the fused element-wise operations represented using -a `linalg.generic` operation. This operation is not tiled in the -`LinalgTileAndFusePass`. So the `ConvertToGPUPass` lowers this operation to -`scf.parallel` loops, which are collapsed into a `scf.parallel` operation with a -single induction variable. This loop is then distributed across workitems using -the `GlobalInvocationId`. The resulting IR is shown below. - -```mlir -func @main_ex_dispatch_0() - attributes { - spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} { - %c50 = constant 50 : index - %c5 = constant 5 : index - %0 = iree.placeholder for "interface buffer" - {binding = @legacy_io::@ret0} : memref<10x15xf32> - %1 = iree.placeholder for "interface buffer" - {binding = @legacy_io::@arg0} : memref<10x15xf32> - %2 = iree.placeholder for "interface buffer" - {binding = @legacy_io::@arg1} : memref<10x15xf32> - %3 = iree.placeholder for "interface buffer" - {binding = @legacy_io::@arg2} : memref<15xf32> - %4 = "gpu.block_id"() {dimension = "x"} : () -> index - %5 = "gpu.block_dim"() {dimension = "x"} : () -> index - %6 = "gpu.thread_id"() {dimension = "x"} : () -> index - %7 = muli %4, %5 : index - %8 = addi %7, %6 : index - %9 = cmpi "slt", %8, %c50 : index - scf.if %9 { - %10 = divi_signed %8, %c5 : index - %11 = remi_signed %8, %c5 : index - %12 = load %1[%10, %11] : memref<10x15xf32> - %13 = load %2[%10, %11] : memref<10x15xf32> - %14 = load %3[%11] : memref<15xf32> - %15 = addf %12, %13 : f32 - %16 = mulf %15, %14 : f32 - store %16, %0[%10, %11] : memref<10x15xf32> - } - return -} -``` - - Snippet 10: Distributing the iterations for pointwise -operations for GPU execution. - -### Lowering to SPIR-V dialect - -The last step is to take the result of the previous pass and lowering it to -SPIR-V dialect. Since SPIR-V dialect is *closed*, i.e. it has a separate type -system, its best to lower all the operations to SPIR-V in one step. This is done -by applying all the patterns that lower all the different IR constructs into -SPIR-V within the [`ConvertToSPIRVPass`][ConvertToSPIRV]. These are - -- [GPU dialect to SPIR-V conversion][GPUToSPIRV]. -- [SCF dialect to SPIR-V conversion][SCFToSPIRV]. -- [Standard dialect to SPIR-V conversion][StandardToSPIRV]. -- Patterns that lower the `iree.placeholder` instruction into a SPIR-V. - -Once applied the resulting IR is in SPIR-V dialect that can be serialized to a -SPIR-V binary. - -[ConvertToGPU]: https://github.com/openxla/iree/blob/main/iree/compiler/Conversion/LinalgToSPIRV/ConvertToGPUPass.cpp -[ConvertToSPIRV]: https://github.com/openxla/iree/blob/main/iree/compiler/Conversion/LinalgToSPIRV/ConvertToSPIRVPass.cpp -[DotAfterAll]: https://gist.github.com/MaheshRavishankar/9e2d406296f469515c4a79bf1e7eef44 -[GPUToSPIRV]: https://github.com/llvm/llvm-project/blob/master/mlir/include/mlir/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.h -[HLOToLinalgPass]: https://github.com/tensorflow/tensorflow/blob/75c40f6bff2faa3d90a375dfa4025b2e6e2d7a3d/tensorflow/compiler/mlir/xla/transforms/passes.h#L67 -[LinalgDialect]: https://mlir.llvm.org/docs/Dialects/Linalg/ -[LinalgFusionOnBuffers]: https://github.com/llvm/llvm-project/blob/ef868a848e6def288d2df7a1b3ebe09463afc8d0/mlir/include/mlir/Dialect/Linalg/Utils/Utils.h#L86 -[LinalgFusionOfTensorOps]: https://github.com/llvm/llvm-project/blob/80cb25cbd555f9634836b766c86aead435b60eaa/mlir/include/mlir/Dialect/Linalg/Passes.td#L30 -[LinalgPromotionPatterns]: https://github.com/llvm/llvm-project/blob/303a7f7a26e2aae1cb85f49dccbc0b5d14e0b2e0/mlir/include/mlir/Dialect/Linalg/Transforms/Transforms.h#L358 -[LinalgRationale]: https://mlir.llvm.org/docs/Rationale/RationaleLinalgDialect/ -[LinalgTileAndFuse]: https://github.com/openxla/iree/blob/main/iree/compiler/Conversion/LinalgToSPIRV/LinalgTileAndFusePass.cpp -[LinalgTiling]: https://mlir.llvm.org/docs/Dialects/Linalg/#set-of-key-transformationsa-namekey_transformationsa -[LinalgTilingPatterns]: https://github.com/llvm/llvm-project/blob/master/mlir/include/mlir/Dialect/Linalg/Transforms/Transforms.h -[NVVMAddressSpace]: https://docs.nvidia.com/cuda/nvvm-ir-spec/index.html#address-space -[PwAfterAll]: https://gist.github.com/MaheshRavishankar/02cdd22f7c99e568f933244b5a679510 -[SCFToSPIRV]: https://github.com/llvm/llvm-project/blob/master/mlir/include/mlir/Conversion/SCFToSPIRV/SCFToSPIRV.h -[SpirvSerialization]: https://mlir.llvm.org/docs/Dialects/SPIR-V/#serialization-and-deserialization -[StandardToSPIRV]: https://github.com/llvm/llvm-project/blob/master/mlir/include/mlir/Conversion/StandardToSPIRV/ConvertStandardToSPIRV.h diff --git a/docs/website/docs/developers/design_docs/cuda_backend.md b/docs/website/docs/developers/design_docs/cuda_backend.md index c0bffeb4f322..a83ddd88f6a4 100644 --- a/docs/website/docs/developers/design_docs/cuda_backend.md +++ b/docs/website/docs/developers/design_docs/cuda_backend.md @@ -1,6 +1,10 @@ -# IREE CUDA backend +# CUDA backend -This document is intended to provide an overview of the design choices made to support CUDA within IREE. It describes both the HAL runtime and the NVVM codegen side. +!!! note - "Authored March, 2021" + +This document is intended to provide an overview of the design choices made to +support CUDA within IREE. It describes both the HAL runtime and the NVVM +codegen side. ## CUDA HAL Driver @@ -9,60 +13,99 @@ written in C following the standards of the rest of the HAL module. ### CUDA library dependency -IREE calls directly into [`CUDA driver API`][cuda-driver]. CUDA library is loaded dynamically and cuda.h header from CUDA SDK is part of IREE third_party project. Therefore IREE doesn't require CUDA SDK to be installed when building iree tools. -At runtime HAL CUDA driver will load libcuda.so/nvcuda.dll library and load a subset of the cuda driver API used in HAL. The list of functions being used are in the file [`iree/hal/drivers/cuda/dynamic_symbols_tables.h`][cuda-symbols] +IREE calls directly into [`CUDA driver API`][cuda-driver]. CUDA library is +loaded dynamically and cuda.h header from CUDA SDK is part of IREE third_party +project. Therefore IREE doesn't require CUDA SDK to be installed when building +iree tools. + +At runtime HAL CUDA driver will load libcuda.so/nvcuda.dll library and load a +subset of the cuda driver API used in HAL. The list of functions being used are +in the file [`iree/hal/drivers/cuda/dynamic_symbols_tables.h`][cuda-symbols] ### Driver -There is no direct equivalent in CUDA to the HAL driver abstraction. We use it to hold the symbols loaded for all the devices. +There is no direct equivalent in CUDA to the HAL driver abstraction. We use it +to hold the symbols loaded for all the devices. ### Device -The equivalent to HAL device in CUDA is the `CUcontext`, it holds all the state related to memory allocations. +The equivalent to HAL device in CUDA is the `CUcontext`, it holds all the state +related to memory allocations. ### Command buffer -We implement command buffers using [`CUDA Graph API`][cuda-graph]. Using the Graph API allows to easily encode fine grain dependencies between dispatch without having to create multiple streams. -Note that Graph API is meant to be used for command buffers that can be recorded once and used several times and there may be a performance penalty to using Graph API for direct command buffer. It is likely that we will also have a pure stream implementation in the future if we see performance problems with direct command buffer usages. +We implement command buffers using [`CUDA Graph API`][cuda-graph]. Using the +Graph API allows to easily encode fine grain dependencies between dispatch +without having to create multiple streams. + +Note that Graph API is meant to be used for command buffers that can be +recorded once and used several times and there may be a performance penalty to +using Graph API for direct command buffer. It is likely that we will also have +a pure stream implementation in the future if we see performance problems with +direct command buffer usages. ### Event and Barrier -In HAL Event and Barrier are used for GPU<->GPU synchronization either within a command buffer (Event and Barrier) or between command buffers. +In HAL Event and Barrier are used for GPU<->GPU synchronization either within a +command buffer (Event and Barrier) or between command buffers. -The current implementation ignores events and barriers and serializes all the nodes of the graph in order to have a conservative but correct solution. +The current implementation ignores events and barriers and serializes all the +nodes of the graph in order to have a conservative but correct solution. -The design we plan for the future is to map dependencies within a command buffer to graph dependencies in the CUDA Graph API. When an event is signaled all the leaf nodes of the graph will be saved in HAL data structure and when the same command buffer waits on the signal we will add all the nodes as dependency to the future nodes added to the graph. +The design we plan for the future is to map dependencies within a command +buffer to graph dependencies in the CUDA Graph API. When an event is signaled +all the leaf nodes of the graph will be saved in HAL data structure and when +the same command buffer waits on the signal we will add all the nodes as +dependency to the future nodes added to the graph. -For simplicity we always serialize command buffers sent to the same command queue. +For simplicity we always serialize command buffers sent to the same command +queue. ### Allocator -The allocator will forward allocation requests to `cuMemHostAlloc` for host accessible memory and `cuMemAlloc` for device only memory. +The allocator will forward allocation requests to `cuMemHostAlloc` for host +accessible memory and `cuMemAlloc` for device only memory. ### Buffer -CUDA buffers are represented either as a host pointer or a device pointer of type `CUdeviceptr`. +CUDA buffers are represented either as a host pointer or a device pointer of +type `CUdeviceptr`. ### Executable -HAL executable maps naturally to a PTX module. The compiler will generate a flat buffer containing a PTX text module as well as a list of entry point function names and the workgroup size associated with those entry points. +HAL executable maps naturally to a PTX module. The compiler will generate a +flat buffer containing a PTX text module as well as a list of entry point +function names and the workgroup size associated with those entry points. ### Semaphore -Timeline semaphore is used in IREE to handle coarse grain synchronization for CPU<->GPU, GPU<->GPU and CPU<->CPU. The interface follows closely [`Vulkan timeline semaphore spec`][vulkan-semaphore]. -There is currently no simple way to implement this on CUDA. There are several solutions discussed on this [`IREE issue`][semaphore-issue] but no obvious solution. For now we force CPU and GPU to be synchronized after every submit to ensure correctness and ignore the semaphore. +Timeline semaphore is used in IREE to handle coarse grain synchronization for +CPU<->GPU, GPU<->GPU and CPU<->CPU. The interface follows closely +[`Vulkan timeline semaphore spec`][vulkan-semaphore]. + +There is currently no simple way to implement this on CUDA. There are several +solutions discussed on this [`IREE issue`][semaphore-issue] but no obvious +solution. For now we force CPU and GPU to be synchronized after every submit to +ensure correctness and ignore the semaphore. ## NVVM Codegen ### NVVM and PTX -NVVM is a CUDA specific IR composed of LLVM IR and NVVM specific intrinsics. It can be compiled to PTX text using LLVM PTX backend. NVVM has an associated dialect in MLIR that translates 1:1 to NVVM intrinsics. This is what we are using to generate the PTX kernel code. +NVVM is a CUDA specific IR composed of LLVM IR and NVVM specific intrinsics. It +can be compiled to PTX text using LLVM PTX backend. NVVM has an associated +dialect in MLIR that translates 1:1 to NVVM intrinsics. This is what we are +using to generate the PTX kernel code. ### IREE flow -IREE's [`target independent codegen`][codegen-passes] converts the compiler input to Linalg on Tensors. Afterward IREE will call the LinalgToLLVMGPU codegen passes. +IREE's target independent codegen converts the compiler input to Linalg on +Tensors. Afterward IREE will call the LinalgToLLVMGPU codegen passes. -Once we get into LinalgToLLVMGPU passes we first do bufferize to generate Linalg on Buffers. Then we apply MLIR generic passes to convert linalg to SCF dialect and then SCF to Standard dialect. After that we convert Standard dialect to LLVM+NVVM dialect. +Once we get into LinalgToLLVMGPU passes we first do bufferize to generate +Linalg on Buffers. Then we apply MLIR generic passes to convert linalg to SCF +dialect and then SCF to Standard dialect. After that we convert Standard +dialect to LLVM+NVVM dialect. ## Example @@ -108,4 +151,3 @@ EXEC @add [cuda-graph]: https://developer.nvidia.com/blog/cuda-graphs/ [vulkan-semaphore]: https://www.khronos.org/blog/vulkan-timeline-semaphores [semaphore-issue]: https://github.com/openxla/iree/issues/4727 -[codegen-passes]: https://github.com/openxla/iree/blob/main/docs/design_docs/codegen_passes.md diff --git a/docs/website/docs/developers/design_roadmap.md b/docs/website/docs/developers/design_docs/design_roadmap.md similarity index 99% rename from docs/website/docs/developers/design_roadmap.md rename to docs/website/docs/developers/design_docs/design_roadmap.md index de61593528d3..309e5f909a45 100644 --- a/docs/website/docs/developers/design_roadmap.md +++ b/docs/website/docs/developers/design_docs/design_roadmap.md @@ -52,7 +52,7 @@ one place where IREE will warn about performance issues, allowing programs that perform suboptimally but encouraging authors to adjust their input model to enable better behavior. The IREE VM also has specific support for hiding readback latency in an efficient way via -[coroutines](coroutines-for-batching-and-cooperative-scheduling). +[coroutines](#coroutines-for-batching-and-cooperative-scheduling). The most common case we are currently seeing in the IR is that of dynamic copies where the offsets are dependent on the result of previous computations. Source @@ -422,7 +422,7 @@ not be possible to continue recording commands even if we are able to ensure execution is appropriately synchronized. This is where indirect dispatch, [predication](#predication-of-flowdispatch), [indirect command buffers](#indirect-command-bufferon-accelerator-execution), -and [VM coroutines](coroutines-for-batching-and-cooperative-scheduling) can all +and [VM coroutines](#coroutines-for-batching-and-cooperative-scheduling) can all help cover for the times where we are unable to transform away the indirection or emit shape logic without data dependencies. diff --git a/docs/website/docs/developers/design_docs/dynamic_shapes.md b/docs/website/docs/developers/design_docs/dynamic_shapes.md deleted file mode 100644 index cc59dc409820..000000000000 --- a/docs/website/docs/developers/design_docs/dynamic_shapes.md +++ /dev/null @@ -1,166 +0,0 @@ -# Dyanmic Shapes - -NOTE: Effort is being made to make this facility generic so that it can be -eventually upstreamed to MLIR in some fashion. However, because MLIR lacks a set -of frontend ops and generally does not currently have any frontend oriented -infrastructure, it is being prototyped within IREE in order to find a working -set of ops and algorithms. - -## Levels of dynamicism - -In general, there are three levels of shape information that can be present in -the input IR (or trivially derived by applying some form of shape inferencing -algorithm). Each additional one imposes more work on the compiler and runtime, -so generally, the implementation progresses by addressing each once the former -is well established: - -1. Fully static shapes: No tensors have dynamic dimensions. All tensors are - ranked. -2. Ranked Dynamicism: All tensors have ranks, but some dimensions may be - unspecified. -3. Unranked Dynamicism: Some tensors have indeterminate ranks. - -At this stage, *Dynamic Shapes* in IREE refers to supporting dynamic ranked -dynamic tensors, where some dimensions are left unspecified at public function -boundaries. It is expected that once this is solid, some support can be -considered for unranked dynamicism, and it is further expected that will entail -new ops, algorithms and runtime support, apart from what is needed for ranked -dynamicism. - -Within the category of Ranked Dynamicism, it is well known that some dynamic -dimensions are easier to deal with than others: in common DNN use, outer -dimensions are much easier and more common with respect to code generation and -kernel fanout than dynamic inner dimensions. - -While the shape handling machinery is relatively generic, we expect that real -backends will be limited with respect to how much they support all combinations -of dynamic dimensions. Eventually, IREE intends to solve this by having -relatively robust CPU fallback for fully dynamic cases and actionable warnings -that pinpoint when more specificity could increase performance. - -## Compiler Frontend - -In general, the IREE compiler frontend should accept modules containing -functions with operands/results that have dynamic dimensions. Such functions may -also have runtime dependent shapes in the form of `GetShape`-style ops which get -a shape from an arbitrary tensor, perform some arithmetic on it and use the -results elsewhere. - -### Shape dialect and lowering - -IREE is introducing a `shape` dialect with a handful of ops and transformations -that are useful for materializing dynamic shape computations in terms of high -level ops on tensors. - -#### Types: - -* `ranked_shape`: This value type represents the dynamic dimensions of a - partially known, ranked shape. It is used early in the compiler to represent - anywhere that dynamic dimensions need to be passed (i.e. function - args/results, etc). At lower levels of the compiler, it will generally be - dis-aggregated into loose SSA values. This type also carries the datatype - used to represent the dimensions. This is currently fixed to i32 but may be - leveraged eventually to use smaller integer when such things are known to be - legal. - -#### Ops: - -* `get_ranked_shape`: Takes a tensor SSA value and returns a corresponding - `ranked_shape`. Early in the compilation flow, anything that needs a ranked - shape should add such ops so that the compiler can later determine which - shape computations to materialize. Getting the `ranked_shape` of a static - tensor should yield a constant. -* `tie_shape`: Takes tensor and ranked_shape SSA values and returns the - tensor. This is used as a junction point by the shape materialization passes - to know at various points precisely what the shape is. -* ... TODO: need `get_shape_dim` and conversions to/from 1D tensors and loose - SSA values. - -### Materialization - -#### Function signature expansion - -Early in the process, all functions should have their arguments and results -expanded so that any dynamic tensors in their signature will gain a new -argument/result for the corresponding `ranked_shape`. This is done by expanding -the signatures and for arguments, inserting placeholder `tie_shape` ops which -preserve the association for later materialization. For results, -`get_ranked_shape` ops are inserted. - -This is carried out by the `iree-shape-expand-function-dynamic-dims` pass, which -uses the conversion framework under the hood to perform type expansion. - -This pass is typically done early in the compiler flow. - -#### Shape dependent codegen - -A lot of scheduling logic will need to access shapes (i.e. allocation, workgroup -size calculation, etc). In general, this should all be done based on a -`get_ranked_shape` op and corresponding `get_shape_dim` ops. For fully static -cases, these should reduce down to constants. For dynamic dimensions, the -`get_ranked_shape` ops serve as anchors where later parts of the compiler know -they need to materialize shape values. - -#### Materializing shape computations - -TODO: We have a sketch of this but are still proving it out. - -Generally, it should be possible, for any `get_ranked_shape` op, to trace up the -use-def chain and materialize shape manipulation arithmetic. Once materialized, -a `tie_shape` op should be inserted to memorialize the junction. Eventually, -every `get_ranked_shape` op should be follow a `tie_shape` op, and the -canonicalization rules will elide the `get_ranked_shape`. There is complexity -around blocks, control flow, etc, but this basic algorithm should be workable. - -Work is ongoing upstream to provide a facility to register shape functions with -ops, which would provide a dynamic, dialect independent way to know what -arithmetic to materialize. However, in most cases this is not necessary. The -built-in traits around types and sizes will allow most propagation to happen -without shape functions. We intend to start with a static set of cases for the -rest in order to prove the concept. - -#### Scalarization - -TODO: We have a sketch of this but are still proving it out. - -It is quite common in real-world DNN usage to get the 1D tensor representing a -shape and perform arbitrary tensor ops on it (usually basic arithmetic, slicing, -concating, tiling, etc). While this is perfectly acceptable from a correctness -standpoint, it is usually not performant: shapes are typically very small one -dimensional vectors, and computations on them are usually trivial to reduce to -small sequences of scalar machine code of a form that CPUs are very good at -executing. Further, we often want to run these calculations eagerly when -dispatching functions, etc (i.e. to pre-allocate buffers) and having them -isolated (versus treating them as arbitrary dense computations) can be quite -valuable. - -We expect that the type bracketing that happens with `ranked_shape` and the -corresponding ops will make it easy to write some simple DRR patterns to -identify such shape manipulation sequences and lower them directly to regions of -`vm` ops operating on scalars. Such regions can be retained and directly emitted -when lowering to the `vm` dialect and/or CPU code generation and would run with -low overhead along with any other scheduling code. - -While an optimization, we suspect this is an important one. - -### Shape inference - -TODO: This is mostly placeholder - -There is work happening upstream to implement MLIR-integrated shape inference. -In the mean-time, IREE expects that the input to the compiler has already had -some shape inference performed on it. In practice, for TensorFlow, there is a -pass which applies TensorFlow's pre-MLIR shape inference mechanisms to derive -such things. This has limitations but is a reasonable starting point. - -## Compiler Backends - -TODO: This is mostly placeholder. - -Much of the newer structured-ops based codegen is capable of working (within -bounds) with ranked dynamic shapes without much work. Given the lack of an e2e -story, much of this has been done "by way of code review" and there are -certainly issues to be resolved. - -In addition, there are several ABI issues and negotiations with the backend that -still need to be fleshed out. diff --git a/docs/website/docs/developers/design_docs/function_abi.md b/docs/website/docs/developers/design_docs/function_abi.md index baa115451cab..50c225d61a23 100644 --- a/docs/website/docs/developers/design_docs/function_abi.md +++ b/docs/website/docs/developers/design_docs/function_abi.md @@ -1,4 +1,10 @@ -# Function Signatures +# Function ABI + +!!! note + + Authored December, 2019 + + Updated August, 2021 A key job of the IREE compiler and runtime is capturing function call semantics from the originating system and providing mechanisms so that invocations can be diff --git a/docs/website/docs/developers/design_docs/hal_driver_features.md b/docs/website/docs/developers/design_docs/hal_driver_features.md deleted file mode 100644 index 2a423e9f4b6d..000000000000 --- a/docs/website/docs/developers/design_docs/hal_driver_features.md +++ /dev/null @@ -1,120 +0,0 @@ -# HAL Driver Features - -Heterogeneity is one of IREE's core design principles. IREE aims to support -various accelerators for compute, ranging from general purpose CPUs, GPUs, to -other special purpose accelerators. IREE provides a -[Hardware Abstraction Layer (HAL)][iree-hal] as a common interface to these -accelerators. IREE exposes it via an [C API][iree-hal-c-api] for programmers and -an MLIR [dialect][iree-hal-dialect] for compilers. - -Heterogeneity inevitably means IREE needs to provide a solution for managing -different features on different accelerators and their availability. This doc -describes the designs and mechanisms. - -## General HAL Driver Features - -IREE uses compilers to generate native code for each accelerator, serialize the -native code, and embed the code in one flat byte code following FlatBuffer -encoding format. The native code embedded in the final FlatBuffer file will -indicate the target architecture and required feature sets. At runtime IREE -selects a HAL driver meeting all the requirements to dispatch the workload to. - -[TODO: describe the HAL functionality, C API, and dialect abstraction] - -## Vulkan HAL Driver Features - -Vulkan has many mechanisms for supporting different hardware implementations: -versions, extensions, features, limits. Vulkan uses SPIR-V to express the GPU -program but Vulkan is just one client SPIR-V supports. So SPIR-V has its own -mechanisms for supporting different clients: versions, capabilities, extensions. -The mechanism in these two domains bear lots of similarity, but they are not -exactly the same. We need to bridge these two worlds inside IREE. - -IREE has its own [Vulkan dialect][iree-vulkan-dialect], which defines the Vulkan -target environment, including [versions][iree-vulkan-base-td], -[extensions][iree-vulkan-base-td], [features][iree-vulkan-cap-td]. These -definitions leverage MLIR attribute for storage, parsing/printing, and -validation. For example, we can have the following Vulkan target environment: - -``` -target_env = #vk.target_env< - v1.1, r(120), - [VK_KHR_spirv_1_4, VK_KHR_storage_buffer_storage_class], - { - maxComputeSharedMemorySize = 16384 : i32, - maxComputeWorkGroupInvocations = 1024: i32, - maxComputeWorkGroupSize = dense<[128, 8, 4]>: vector<3xi32>, - subgroupFeatures = 7: i32, - subgroupSize = 64 : i32 - } -> -``` - -The above describes a Vulkan implementation that supports specification version -1.1.120, supports `VK_KHR_spirv_1_4` and `VK_KHR_storage_buffer_storage_classs` -extensions, has a max compute workgroup invocation of 1024, and so on. - -The above bears lots of similarity with the output of the -[`vulkaninfo`][vulkaninfo] utility. That's intended: `vulkaninfo` gives a -detailed dump of a Vulkan implementation by following the structures of all the -registered extensions to the specification. We pick relevant fields from it to -compose the list in the above to drive code generation. These are just different -formats for expressing the Vulkan implementation; one can image having a tool to -directly dump the MLIR attribute form used by IREE from the `vulkaninfo`'s JSON -dump. - -When compiling ML models towards Vulkan, one specifies the target environment as -a `#vk.target_env` attribute assembly via the -[`iree-vulkan-target-env`][iree-vulkan-target-env] command-line option. At the -moment only one target environment is supported; in the future this is expected -to support multiple ones so that one can compile towards different Vulkan -implementations at once and embed all of them in the final FlatBuffer and select -at runtime. Another command-line option, `iree-vulkan-target-triple` is also -available to allow specifying common triples and avoiding the lengthy target -environment assembly string. `iree-vulkan-target-triple` will be overridden by -`iree-vulkan-target-env` if both are given. - -Under the hood, this Vulkan target environment is then converted to the SPIR-V -target environment counterpart to drive code generation. The conversion happens -in one of Vulkan dialect's [utility function][iree-vulkan-target-conv]. The -converted SPIR-V target environment is [attached][iree-spirv-target-attach] to -the dispatch region's module for SPIR-V passes to use. - -SPIR-V's target environment is very similar to the Vulkan target environment in -the above; it lives in upstream MLIR repo and is documented -[here][mlir-spirv-target] and implemented in SPIR-V dialect's -[`SPIRVAttribues.h`][mlir-spirv-attr] and -[`TargetAndABI.td`][mlir-spirv-target-td]. - -[PR #3469][pr-3469], along with patch [D89364][d89364], shows an example of the -changes needed to add support for the -[VK_NV_cooperative_matrix][vk-coop-mat-ext] extension for Vulkan/SPIR-V. The -overall steps are as follows: -1. Add the enum corresponding to the extension to `VK_ExtensionAttr` in - [VulkanBase.td][iree-vulkan-base-td]. -1. Add necessary capability bits to [`VK_CapabilitiesAttr`][iree-vulkan-cap-td]. -1. Add a corresponding attribute to the `SPV_ResourceLimitsAttr` in - [TargetAndABI.td][mlir-spirv-target-td]. (Note: The corresponding SPIR-V - extension is likely already defined in - [`SPV_ExtensionAttr`][mlir-spirv-extensions-attr]) -1. Convert the capability bits specified in the attribute added to - `VK_CapabilitiesAttr` to the attribute added to `SPV_ResourceLimitsAttr`. - -[d89364]: https://reviews.llvm.org/D89364 -[iree-hal]: https://github.com/openxla/iree/tree/main/iree/hal -[iree-hal-c-api]: https://github.com/openxla/iree/blob/main/iree/hal/api.h -[iree-hal-dialect]: https://github.com/openxla/iree/tree/main/iree/compiler/Dialect/HAL -[iree-vulkan-dialect]: https://github.com/openxla/iree/tree/main/iree/compiler/Dialect/Vulkan -[iree-vulkan-base-td]: https://github.com/openxla/iree/blob/main/iree/compiler/Dialect/Vulkan/IR/VulkanBase.td -[iree-vulkan-cap-td]: https://github.com/openxla/iree/blob/main/iree/compiler/Dialect/Vulkan/IR/VulkanAttributes.td -[iree-vulkan-target-env]: https://github.com/openxla/iree/blob/b4739d704de15029cd671e53e7d7e743f4ca2e35/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/VulkanSPIRVTarget.cpp#L66-L70 -[iree-vulkan-target-triple]: https://github.com/openxla/iree/blob/main/iree/compiler/Dialect/Vulkan/Utils/TargetEnvUtils.cpp -[iree-vulkan-target-conv]: https://github.com/openxla/iree/blob/b4739d704de15029cd671e53e7d7e743f4ca2e35/iree/compiler/Dialect/Vulkan/Utils/TargetEnvUtils.h#L29-L42 -[iree-spirv-target-attach]: https://github.com/openxla/iree/blob/b4739d704de15029cd671e53e7d7e743f4ca2e35/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/VulkanSPIRVTarget.cpp#L228-L240 -[mlir-spirv-extensions-attr]: https://github.com/llvm/llvm-project/blob/076305568cd6c7c02ceb9cfc35e1543153406d19/mlir/include/mlir/Dialect/SPIRV/SPIRVBase.td#L314 -[mlir-spirv-target]: https://mlir.llvm.org/docs/Dialects/SPIR-V/#target-environment -[mlir-spirv-attr]: https://github.com/llvm/llvm-project/blob/076305568cd6c7c02ceb9cfc35e1543153406d19/mlir/include/mlir/Dialect/SPIRV/SPIRVAttributes.h -[mlir-spirv-target-td]: https://github.com/llvm/llvm-project/blob/076305568cd6c7c02ceb9cfc35e1543153406d19/mlir/include/mlir/Dialect/SPIRV/TargetAndABI.td -[pr-3469]: https://github.com/openxla/iree/pull/3469 -[vk-coop-mat-ext]: khronos.org/registry/vulkan/specs/1.2-extensions/man/html/VK_NV_cooperative_matrix.html -[vulkaninfo]: https://vulkan.lunarg.com/doc/view/latest/linux/vulkaninfo.html diff --git a/docs/website/docs/developers/design_docs/hlo_to_linalg.png b/docs/website/docs/developers/design_docs/hlo_to_linalg.png deleted file mode 100644 index 469ed26960d2a72055e3f56473cc970e62b1444f..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 51823 zcmXt91z1$w(_cjdLFteN0bxN>y1S9?7HR3O1*Ji{W9d%m5|Hlh?(VJyzKieozt6+N z-Q9c7nKLtI=FD$qg5 zr9*_fz=u~Rg0g}jP^K%6PYpeD3Fpj`@iJ9Fv<6{_<1R-zrjaX^9|ECPRr(sAnsExRLX%Cx3M+ba6HjI;*qAzCFMh`uKQn z{RW^0xV#G_ZoOXocQ5l&5bOW%kbt}Nng#B()GFgCZ?Qi^S#cgs&Vv+u~)Yl@@)=2*5L@CukxsK;{RJqpa&x^nA)a%adEp zY!=sAK3ZAm&m06Hh>LmLIEnr~gyc=rh~& z<;jsi@}IBA;xEpBQE*OJ>7JEWM>piNwyJ0^pWco1o^B}I=NDIoyk=j`dlmxAw|sPX zP<~`E@@%iRWj=H|;GC*3PhM!wXlDlZuwHSgla`%khXr`j+aKk5tzt=UbsZut0=0e* zx9A!17dRcx*R;ve0}uR1rlYv+(|5PNaM$c_?DiKCY&X5nbgiu9BZAVuP$PxsbubVe zpJ={p^0+u(!JRt^e}0q2VW$}Q-AQMW$zo#TRVhwPD@W$@5BUGqQ-qb{yoq|l++A8$ zY-Z3(2k(XZ{Nq2i#b;hPEqO=O9Y?PVE!5xp1$4JEasUtaVsq!Ga1q4G+&jKY4d%1f?LxmE`W#+3b)RVyONrM1Y74Qjw z`vCH8pgcSa_ZlJ`h`oBc?&)OadI!9gE+~u!aL&S#NIPJpEp(m|yXu*!k#w@>@v4KP z8f$MqJb&|Ve3{wX8?#Uj;kfMOjXX2?4H8MNHtYlc#tSP+s%IitOc8+vvyLiHHN_LD zaXcQjJUQ!hgXlJ=ogvL%y~x0`gM_@TJTzkbL1OaX=EqV?9Da4*?WRZuA#cT08ReTyK{vkc{hLa9LlbP8C+%TU|1?J~Fz z*XK$L!fQWuA+7r)o|@~$B+ug@{iFu|0apUnm$7FR1HO}d=E;9|8D{@)Q_Jo= zmx5oA&JYjDjjh3Cug6;| zFY*4emdC#M0mY}s2g~o@Z@4uHiHL5aB>Lm%xSfv%mKCCUzYQj{Zyy8n_8!iVX*6j@ z#dgN}{G{O~dr<%ALux6{V)`{RMc&mO&uj`7iH?nj)> zBG1dCP=5>)87op_Xr_p)vX~y4nu==y_$TIe6G0&-sh#p|B)tXqlgq}|RlW{4{L_1H z*G7Bsq?CmZ6xm>KcZbBwmxbf%xBv-zgm0oLv?}Zb1zUX(P}Mb3nhT-DEI*ZsR7)Bh z_q&!A_U5XKoSH}yKScN5aI1>^sqQ>lsJGvrqI(RB5(AYw{iX_AyX z0E2nlrBkr8$MgJ%ExH)^Jd7A5n&qURjcih~JNo^Nd-O;?wApS`%FGN*i;IN?1I8-W zrnEK3HrM8M*cmC+;a2=h{~A<8HQ^^F7B`&O{uJ?2J+q=hOge5R;uAhPnMTd8G*i-7 z3ZN0$e05Br+>uXHe<#R>;@VU~gqBE0YQ|}T7{(_xURbSci}w@>VHcgOwo)O0KFt-X zL~Zs*3bjfhYB((O1;yj?kD>va?;UO{YneZ!{F zeAioaps0J89gIo_N+>zsru4j*>iVVdeZi&7)t)fGhqcaLoXbcmxhs?)^3z;1FMqid zFzA#_e)8AFt_oooO&J{4-7b0|lg7GwRidvOna3}`-+eS{MKB6v$Fqazfv+@tT_E`u zAt4F{g($;rse{8&R(kdDi3uDCWRTrD?&kJXsnC}cN&j@lph3R8H*A!)g)iFigj_o~ zDCG@~EO`M&apq)%^@~OcY&>j`j1MdWa%a{%kQmN?e?ZLNiv*TBhj1n&$5*dr?Yzx6 z=MhgksqLP-#=r9#f@Z@gSVaqdpg<&HM4mi9vrJIz-M&m)#B6oQTq6!fNG!CC*S_Yg zU89nwz~*FX=MhpNBHmfYA>?&dXAeLpKM39VtQQ&2VQFf^1$Gr7ncuP~J1;Sl*Ov9X zS7{NktJ)&+JdoDAH-NAk#-C}?I~$XZB7wN9?A-;X zeKV3`$QGl!4B5&|v|>&>VZga1&yRWd80x$E9zgMPI&fcyGsYlcoWr78_`@BItS1utzcf*##M z%X~C(SU(10FsEIDEj-+6F}`YB!Yc@E;iT5F_K1F_*}tM@IMoyOQF@yy?eh3{CFan8 zoa-dIM9A2^H<$3f^w*H0J8jbA9378kZj}BtzJIK_5vU-70FDW0fi>_?SAU z`7V~PH=HQ3M#MiJH$%(&YC0;seX9a(o!0Qxny~x|tljS@qi(+;D>BE2|1PKsR>~J- z-ROk!=ev@#wph0X3mwT=eYf3S)JPwDg2(rHFTh#mm z7x0y)&b%W%C=ew6IFd zp9wDnV-%Zyx{XgI!7GTVWYL}newp2xpZ}fUvTrkkb-WYpuWW~OXlK+4PBb9@Rfj(* zRAs$vb^2BE8--V`9k6Sw!wHjB6FECAVKxy+`n&1@k}8l>KZm(;L-7gdOBMl&z}N69 zz5($k*Cm0>$12#X%D5`caS;TZrAnBPh73XF7?(-i47Q*7A_{gKZIkXq@gkv!!T!UP zBC17uxe{*$vHm*DjBa;^;A$)e4ux=PY>8Czdi!_%%|#pLy(Y9k__6=xMX@d%Vy4O# z@9GE1%MMc)v{-GqcZA3BF=dZP)fggvR0I|H9}4GIqzD#}f}=I$KWGR08!RzVgo8SM-~RnF82IA)8gf2=S{6W9OU*Eluw*t3BO+} z(`e~u6Scpam)2GG&U>?Bi?k4*IQmM3#c*H;UDL2iITK&jezK6nf-I0G7)5lz)gC%J z9ccHq?nY^rej2|Ki~;y;<$xOjMg|5sZl~bn;cCly92F=?-dVlVOi=K}t?1b7fu6R7 zxrs@D9t9{we^5}Fu7`=acEx!c#f;&Q>*6V;2tQ$bgp_74xfX|;mS&LKL;pA~$0Oj6 z*kw|GQ6aZ>(Q$5|AR0U);=ba}-eeIM>Y4A{7Yy!RLE_TASKfCjMl6&qnpam>fNMSg zluh&eWdhYcN~UinH8E&FK*Me*SYRSFq*F}FR3RY zTF4zx0A2ALD-l-sf^ocg?IBypX7fH>O z6Dx{jMq_Z(A*{bhX%v&3Z}hFoH=L&Y!!uQ6o!*4QgPHop^Zn5Q3&nLO;(@_7t|AO-pb-13Y8S zYg~~|EK$xUUJsY+ybMY00I{TL1AH_QGu0~VvPRpo)9Jvd%h^UtVZ@ReN8~V+#M|JZ zrPvFE*`w86ai5e}xD1Tov6Oe?ym2O`ufJ?MmQ2=NOpK8R3JH^0S%w&X0(tl0oh1kD zOTXAt_Cd+y$*3i}tTHw=5xrBcto)#WAk9Uq8H$v-vFQK=JZ!VoM^?e_Un%@H-k{za z%aPu=kzD^tO(6#YWd+)ssj9}+VfTI>=2<^Dkg;;VItcnm0jlmX_sM4#!;h)Rtt3{GmE_DapNq?lDn9;@F1Gl69pzI{1W6lRUhm8^_wM{NNZt)a1LG;L9W~C!?a+8K>h~(x;;-Sh|7j^57R&m()l0c-FfLg|Lew-oF9gey7q@zTT-u#_IDCu{k z;~hgi7u;fnxXHWGYE@tcrXHq2z~ZsaEzHf$q2Muhw6?yOOa+0A10NhLv^8j84&*T} zmD`fIX))Aj&UhCJL;c$vlkWFd7u!}c@v*}%i9{?af6y<>CsG`iZR=?>f$zQv@BG`@ zmARiiD#rIb@2@&%Xh0ww-|Ce~(*Yr&-JVGDQ4d^!WL5+PA(eRz>_lq^DPL=A1(u>G z(!$gniwOa>+0Va=lILT8y><2Z|G~%!3$k-^PNPLuF~xz=R0T|^jrcEL(U*Zpi+KVu za$Ez}rvDTaUrtDFVJ5!c+8&ei|d&ykl2Uwj-fiR&S zbV(cl;zqY4hsM)PT$u0gV%>qUBF>-Add}^JwJ4|nBY&=i6oRx!a(>WPesii(8=?NC~NUS8hOGaL}8@P|rt6hEy{ zGX9j!^z1BN5LPkjPil4cR|@%_bD`FLaPZ2Immq->4)?3?U%rIZt(NBH<%NVG(^JA1 z{UAk}0>Ucucl~JysKgVdT&&+`7_|q0xX4lDehf8AVsT+XUK}Y7z$X699^>?Qu-x*L zuaGrz157O=Ngq$b>!vJ`ZCZ`|zlW4D&!~X#G!F)Y>3KbuosMS2V7%1jZWo_8*-BgOole-;Dpb1>9J5~ll}IP#dN9X z^&$>nXubkGqdw%{?nwinn8z9L1mKI2WL-YWB%A=Y`8mj2gl9zP?Cj)+owNhNnY2{h z#%#64bbh|Xe{ajfpH<$B%kjU7Ur7If2?D88U*U!l?jy@Od3t!r3nE<}9n4nVUmxSn zB{7BIGC0}d0TF>JaG<#0vR(77dvoOb26JQ+;Ewu>QxXm>t@(I8i0ejAr0EJ=t~g)t zGC(IV-Oareu$a(azeG$#pJgl<-Sx&$TLE}YRdW3pO67Oepf6=t&%?qJQ{9QJ^~dEj zUC;+mg$6)#`erLlJ~K0qjg3jn#$CA`A035#tOHitslLqE=ydp1EuH6j!5;IYi~u2g zG@yxkHQ({P0=`aWWPAT7${St6sHiC4B7xpzeOz2za74!4*whL8?`0jSH0F@;M14tb{mfa2sA)(A83Lik*hV_Tr{+px5F(vnIp71oB#kp2h=SOBg!ZDW|yqIiFIm(!>Z>!KLBS~Nw-k7lIOWH>m0 zN??3p&O5h$3t85l08JGc|NVQ3$KojdoX5w<{izZHUN;s+0B-SSQ{}#Uyg%kYANXu$ zW_H|se*|#krW<6w0`TA?IEL_7iB6;H0vfK*#)};vQuBhD>hui;>X3B6fQY{W&>hBC zy{}0;FCtr7_`%)X-Ro@%`57fzwUEc#?G2YtEG$@tn4tHFf)gJe0CUq;onDNC<%tep zQSXf|c19_}+}zwMkj0%3=TuZxyZB%0Ge!w%ln4k*XFLZH1#b-2iRM4Aw_{vyVS#VN zt7C9CVF?(R`5sR9S7ASx!2DZSvFG9YZFC=)zVq<6BmZD};!IIIQoh>w@ z4=lxJ7Zz366XI&6WW3~hQbs4oQI5o17fQoO(CjeAvPQOY#dFPxElziAW6i|*$j|@O z&(QMCry>F|*g_B>)O-inIcc6(8n(8!00IQ}@}+zhy#Hr5Kh1vDOiI^+wC2NFdY_h} z59dKLZ_n*_Ct_OUz`@MyzD<=pyKlYpLbK#Ls0{3NxXq@*b~Stl*tv`f#=VwD>_hXH zu*vNk;Kd;gF26i!Ln&3xbU&zarzQ|J=zQ2civG>*q>cOdA}kX5gP=o#xM{P<+PTWz zFX>}fuWUoHqY}}ZtD4ng&5ysx&7pmLeUA^vUU@k=k<51h>;lRcZ6nb*tGx_mU;3+0 zAWwOgI9qQ|C7JPXf9|5m&n$+(t-BSVX}I=UWYL zJ>{srSs^#i)Gg@5i1zzACDES^Nnj2#l~^_a>lKHV(|4pbQyGJB=4 zj!W@swC}%Ao&o7M4gwTggKP>X9`?gyE73%?L=zhg%^M2|Kn1lVf=Ru!wujU7w6zJb zv8x^TRUcn~bR>`-HlpQz$!1ImM#fhtCkh$yp>v1n)DPCXUd~+EkccxYq`t?f+I+b@ zpls`C-GO9;bcq=Dh12OW_v{ut+Dh!uz+<==9N46$LZ>7R$l-SamHqk_I*y=87}^?k<)QEe(YL--+1qCX zLafmuQK?t`o8rYYXrzHj+-JB ztsWg%K^<&7R|Fs5sChJT+4nX~o=u%{m}HgUO4hiFVNl-$gZ;rFQ7$w5OmQu%K1JL@ zYMJUBG35Q~Jw&GOB^3QswaZ`5l+Wox_QuL{gto7~#{-(u4PU}vXZS*=_SX#WEai8} zwx3nY$QreE7d8xTYQn5I+5YsjrYWEGUJMMfSZic6#II6DMdQ`8n`HAZrwsXc{#!}tK69+#yVtbh-dyFU8&tr;( zO#_GR1*`ZRP21Bo1U>(=Rh^~Sh;H42Ewe4JWcSgMp!O^CvU7Zxt0woV1oy-GVT4-! z?Z&Pdr*oZTAw0YY{pFYk1BRB#6IxQyq>7=Sxq~IUS$w3DUZylnqwV#NuoLr~s#@ZN z5Wtz0XxB%n_|oe&6M6!^9h3kH!lBTQNDyojlaN3pC{Ffzi_ccCJ*2n;ILMxI4Ik+Y zibfsIOU%ZR(RO>6aN{b9sez?yg1ME80jK39#oob%w{e;}R486z&U+DjpFKVdmI@i^ z7r@r=rj8ubecbZIc?iUZ&ERe;qSl6K^Xd{C!Y}Z%vB7u+vT1_k^+KZsw6*oWsXm!* z^kg_S;7k-ABQm%*7b`>IgB+c0_i~#Sv+2IijiE-v#-Oh!TH#P+lW599xxGy4Ji!`4L5gki(I5>UwpdYS8!EbT^(7TeEAm9cg+*g2Oxh# zKbo)Bk6&vwAI5Nog&KVPzA-zvt-s2kRH3kvlvS?AGV~_Lm(D;Z^z zqHHuvJi4U~Mg&qoqD036$ABAE&LSB0s4OQ-O=8vJz{GJ2Ig_~+sfR29H!KNNIhNcF z**{vrxTf(h!2MAjC|J7~AOvhh0!Tf`|UwC8PvU zKlXBSfHQzeLQeBIdn1{#QML#e;bA$z_Ils>f11-T$1=CD&}sy;kP2X68ji-lbFxXs z6R6A)Nqv9T)z{JP>hVVzYcsp+&*mf~1~*MJ;HGM?2Qkk{7n!MyW-f zuQ9qVoZ2~VdD;q*mqQ`sKqDb}0#5w!2ngrVfTsQZ{(Z5*ak^9|AIE}YrW{W)1GuWW zBitndDQKh=HEk=->kkvum7nm~k>gh8GLst`$we02fb)#phw--qH+<;?+ z(MnA0>OEH5{5ImPfMx8EBuQ}xS{+`P=pOuJPYOZ#>-`Ef8oF9TEML{89JQxML7=L&H_qax-ksbAC zLT>p=E04B?^ZzlZ7#c~O!3DIdp)(3DP~JzP;k>?(3d}YS;V71|)c+m8P7o=Prx7zn zC~3vEE`ucAChupKQ#Vycf2R>-mSqXoyHq52&Z*ktHPzcBa{vDGAqLRVYNL94xvsHH(%%K>g>-S z%`i%V=G`oDX|Z1sRrzkpQcC&bSW+7s5A8Ba$5cSlX@L+91m!fbXM3)~jtayIY#xFZk1a(|cTU_4=btH%THKfC!h zu}~Q|q>}cnsvgW(i?e%Rl@$AL;;>G2noZca^TKu5RKn^$vvDNBOCQ~?i}PH2eTix! zLxCil8wkE554^jZJ6fdmJo!6P_W7r&XWxIJ82+XZ5vNzG#mktewpeoVejT0uMkmjO zQfOHJ-iv__DEHX;cG^s1Ex6I22H1Nz{&Nb#<(P6sYSI-Ib&ZBy)Y{s_*SqHle2Dr^ zLuuA2(8UD8ny9k~vZx$sN`@BgVcVFTk;jL7gjcU7M3IdCG}TwOySjGXgnYk}iVd3P zzMF|}zN%uN93lT)Ior=r`zt7NAZ_?L4P@&1wZ$3OL68>6!;K`Unu1*TtPNS)pfk9| zco^c{;-s%3@!{cOT$Lg$s(rkyg{>-5PT1)XA9rqjjl$>k!^7SFCPS05zcsJ=t23kF z9@82t9YEmtk-k^A+n4r71;xb&MPRrXi*h#cQ-2MajH#(w>Nh}Q)T>N8QcRRdKw z5dybbmu+!fN|e8Tm5vfJ{Q65Gylsij>SeGE{^-ZtHhmR9LR{Ctx0d&_KxK}->@VF1&wXtZUflSX<+g@J5fQmdOHs$o zpGzHiYs{MJT~0sL9~w@$!ZFeY({jIj&0@}CK4|-S&w6E5_^3ht>ttboM20>*fL;p; z;c;r;?~p@~0wc>{uv_43%ARZJtIubyAJfdwX4zhl`2lfd(kDL*3S+N3uo`JAH|r!% zgPD0?c1a2xTkan%$rcypv>rGDxCo(3Fs((0;!XsacB7kN*=gwm{;0(?1uF50QGUm0 z7Q)q0Xd1V3EL6uCg0XXbF}^jHmO9o{d5X;Qat)dCNI|0!_nproJ)9rFb{IYH93X{nHa8@6 zy`bEYEcRRLD=P?BSJ*H({jY!nLQU$1Njwv2M{@S)a=QK@_n zBIbLT#r^vU?~jtU_f&8vh0rLy5L=#72~y%oy->QtN(4+!SO znl4od&dBBS5OQ%$o*=!_43ZcaeO;lSC27TSxw00)jk(2XsZ@T;`O_G7^*A!OAv|e}LTQuLoote-U|xGJK-wX_mk!EPSkvYupy|Nnvg4 zcaGabL}qq&ArWX+6YHM$b>YpsV5D%!{jPhN=+XAE{Qz~~M+%!fU&!6=qxTE1(`cC& zf><)EE3I46RJ9y-*}BQ13;b!1%R0$ExKt-pnSX@Vwwf>6OvCvA3yMaF8*Kf)kg^(_U2er?i$T@ts{t9^SsXGbWPCbb$Jj-p27UW#D0;B z{<$K9LED7{j3v{))VwxtE&cl4-Mt^K40n}&Z^57k(k0Stj4Ld8S|@d5&v>v8%Dqrx z8qze_%l04$R?nr^sSDxz0{wa|BV+q?bB9;Tr9^E@W3W1bXAz!MS4nyW$D!{n zs*OBur2U<(`9wcvwY@&oy;Rre+8jZV+FPRl0x06kc!uqv&bzzJXl=%VAt(u8B%I%6 z_cqd%+oMYM(9ArX78~T2-7x|g{ROqC7C3#Kz4NyrYp@z4g$o|6Tl7F_ZZ#OYk{vxA zeFrpI*TwX-k0E?oSc@z1EpL*yqxoV(T+5k6252kyPJFz05@K@#H{thl_ucywHFqIN}tc9LW`Icu{IB6-x3VBywW*ZO8FrA4^jJ3lDf*WjOf`( zVjtXh+pQvH|8?doZZ1w_iDHy_#2NOUf{{(H4fNV9_!zA`XJA?jIh-!@oO{ZhHRw`=IZ#n|B;Y(ZPBY2L{8syGbg>b+rMe6 zMU|oZN}@Q(;6bUUD{7`IW^Vp-54p{jkR&cbv$xFGYJ;Fb?qXb!Z)1KpUbub;c0H=^ z@3lBl3%VxHz1v(}k`8S?HRtDQOCD3qYSL$_v0Cg?GzG5fqxXM^{%R|K4&2(Zt7T|5 z?Mvl+F)=f*_7?x*dTS7-^P48TjObC{-zQ-BV1$B`U-E+-67-PQ!F0tLmnA$HI9{^9 zMj_feWasAKf~IwS6G6BSCE;?={a*8<+a3T3NnE&HZeELHd#5mcW`NSfbVP5ftv)H0}Xutm-ul0@%RQRDA8 z36I~0geq>g%st4!b9>edbrIZ$E2HiUgR(p9xbU~ss_3DG@sQtH@*k;vHLMB4Tu1=tsh{Ad+=`i zSq*O=<7I!;t0m}OYxi-4sYY<%pIyy8;kkT}&7|fbWGaga6|NcB;Y{0me8FVJp#C;n z9%QfjU?K8iZ*g4C-s^ptj^Xv+K90L*@UhQqx#8>|NGZ;xliK|GyuwBzyXSs5fazjI z66UeDS01yY4+$i3b<;#SRI0_)WZhRgr_HLPziT_XoAi!aUQm4KHQ{mAl04-*NVAp-5ee52z4pG>qbA zS=iiPVKc78kyFd~kh$~sjo1BJ=yTf*wHLXY{36q+GkuoN?!A4SE1IE1%` zwk2cDIso1L!cp3-JNyrLR6CPS07BVj_}OZ7Tcz}m&GB4449tV)$#g}Y`)DWu5%_Sm zFVpV(!@c!GOl)j9)?CYfSS%Rp?^___1aG(~$|q1esM<5izS?*xxSqmYzHgW8WJk3TvpRAG`i8QRMjnUI_po&jVV5X z5l|?!54U(|x#Ix7W;U?hUDdrP&A^KHn(GL&{U2KU+qx5CJ&O#9#v5E3Oze>5FR-f) z?*18*aP3#0z3^H+x+0aId9h}n%5zjQn0zz{p3u2pUDaGvz*W_?BG}$I87Io9YI7bPSfu>WgJvZJG7{P#FCgyy+ zq$ueZZBd48 zzArAoAgbGEvOIc>a9@2Y4-uZNC1+Zw`=q&EmOS4GCWwDyZ>vR0&FuFS%Ta%LuXiCT zj@=nsCF?%JeHQd(sqbOuHwn{~;mzZe%qFVDSG%NAWwR6A-yjM<$XPGdC%@+NA#MN< zhrxmLznqw0f(F97TbA(8tGb4xC#-HJ%;3yUhZgT)omz_S%ZC>m4~f6g`wZ#JjTck2 z-|i&{opv*Df3E-eiPWwje(Hp$X6TWui!cExpr1M1Nbs|pMOo56*mwrLGOX}*qSj4? zA34jE&|KklwrIa;X@Fdtfj~f_t{Lzrdyxyk*X-4aUMdZND9=L&wG+un8Yf5 zN^_masbmb3SGyMPRfzGlz@oi6ZgmVhpDLS_^iAE(Mz)#`N^n_Q9LDI$^O{BYEsWTY z_~ao|NwfD5;g(Z`b}@&QydX(#?$E03{#Zq5`CYkx2<2llk z{t@D-D@)j(now#2MBDXBmSmMEa(P(vGhr3I8qS0Fj@79CjV`0}C%;ra&!n7@twGwx zDgw}6%6Wj3jT1jYxC6ZIGKXmL%GiXfhWg$e_%{7vCXK?HOoxd&_RO zArD<7uifKr>raQ%-ZQr9lGL<;a=`EwTPnv){N)^SC=Ux?tEDkLI6AtI7kAspxqV7b zbFD9IA0yaizR6M?Ugb?1l}yU!)uuW#P*>?K{>DF6P#MN*Z7}|WXhNq!=~AUuRZ78* zdzhDk{`S5OO;xU$pIILk(OOo}P*fr_gnzu8%_~GM(in%R>Ljl8VKR458OWVu+C%)s$ zW)A{00xAi=(8R4JHZd0NXnW65n2zIQCEDPD+r?IhiXu;)d)6=79ljbv-#N`g<=Ilp z{Vop$XqDqUl^NO7Ky@DjzS2TVKMKobwwsiSXSQd;d^uDWvl(&kxtZZ~rTZ~D_r~kb z@o_i2MzYXC??jugA11kzu6R@asas7Y)0Mu6VO?$*3UTuWa@nMHjczbPr}GYbbBN(| zjYDwUN4YB-R+gZti~Xx_^HUAm<_is$+?sWetvMadk?qrm#sg*--n(x8kqZ+H`l#5>19a!}W^V)7rRe$ymr8b!1rlGCxd@U~eU?R(M$m}*IlO%FY_27Xn zS%tDoSlT;E1*5$Sm6ICgg<=epQj}?9COcKncG|vcu`;Tax^O5~w%wmIYHvREZ^{aw ztlbF@bDZ!Bj>N(zCf3!~{`IUjKtIG^JFry6b;lT)Vn-HAO3=GA;%>Zk*f+doP#Mx# zUIXXoM#9rLScg{=L#Uq}PC~23Gm7m^|1h7U=yI%%*MqB{O0xMRsGmXSR_U60un@=F ztw`RbQ0$M}GP-O1yxNyhg$#u!4Pi^S#%2jEovlTy-|<699nb^lylFeasg11oSD)Mw zYgot-6suFT^CB_UeXYRiKwDu+E)U7!m?`5&(+l$($LO?%GZSon?z?2Pb6>GPcgyG+ ze*9%AU-<*$^q#=}z*;i?t%Y}>pHm98=ubSIo3JG{g6`>UF{Vek-9Oylb{7Emxf=OZ zE`!Yi!%(b>HN~D>k_&MOcec_MVl%8u*OsqxE6Qm&JAg++91N zFc+~9P;gTV`;on;+c}w{Q$@fS?nRd*-h8V05FD%Bch?(%FJe|S#5-e3V-l81J9cTL zL0|A z;D3Jt%2D;}eRHk2VznjSrV$Lfm6ZgKaOVIyfa)vLV@lkbj$&q5zVZV1=D*CsjtJ8z z^4^Yd^uT58RQ`4F7_vKiM@0sVmj|kG7yv1JWK&WtnY9&cZ0{diptEQ$E=-msC3yS} zlZMvY3Uf0^x8&V}^lVSJvmgAj_1R~0*fBg8(KX}tl#`O*SHy(A)~3ur)pA`EFS^#9 zTqrBF4ln?0xwOXRY&aLh>9I1%`jlyKvA!_3DpS|BMZ1!+Hs%W!89yoaRIyZs(L5wh*k}dL;=F#Q>7Pwg%WmA6`rH8WnAJ_4$9M?NBocG#G8~uTid3`7zu9Bv;pv7 zkU^6fy$oI9*B-c_4Il#2sgsro+o^A(DUf4}6I)J=G1Qll+FK%oZhtEhR7i@COiDsy zTZ;7Z90JN9<}GaTKwBA(+q%W%sCFmUDusJ+y1Fry3{2l^EEow4Qk7KLD+dpn9(;Bc zQg}v(?R5__c=&i}g6$Lyh5m>#7$KU4N0ycUSlJ9}Ev;g4+q`p20?Pe!BNo{Yr+V<^=f@mP+=X=RC(`e#A_k$~O*~ z!lHyq4|TH_-|8mr)uedRGu4Xj$F;gSWsV<5k0)K3cKfWUZ!aZ6H7*)@I|dongJd&LW3-OMGjN=DqyIWT+v5 z6dD}-0UkJqBO9T26USD4MrSHrP1%iXv6M7e*Kxc!1V6`n+xx2{wjMLitRcY{a?@3L z0b$qY?%YYKVk%tGSa+{=x?+AGSD=eq2=BLspgaz8)|oOAEMPDT@uqIcdG07wo^#A~ zAb4d!HsFRB#J5&X3ys5**pHUVVA0m=_;sr~Y`#rAX1CcHi!&CE4kstbD#rw!Xk&TZ zpPp=f5a@79x;gy59?6@XbM$b_ueDnQsx$bsOL-HTI|jSEWe@2hV-vccta?0_v>t++ z$XXT^xyv^DoGE(Kiv;wFJ@2g+XioUiiILep5oX!%D6cp=pt?k(ymzC4oDoQ`j}I5| zNpAXeE3L;jGjTp*IS)Uijr&*UjT898~S_X$F%RCV9yAeOpuI^wG zodM+4mlrH1_yOeQ75#k*tj^=-g8c}XZzPCw-^J&eZ2SG3#|r`MTK%VuJc%b!+?VV0 z+l)-2tE;HS94@a4pAsjc(0PXq&;HB|n8SCTWwu8y74qxcWs46J@R~1meTS;GzEH=n z_Ti^vP75SLXg0V^iTWImT=+h`$t(~3N7~0vBe&)oIn8gtLam7>Xro*|Z(8@-@99GLHR|Rm)BscwwL=K;Wu|FvUP( z0FRL4C_QJ=XJY%PU0WCR5cwk2(-PmF&XIGF3l)z9P0~X65Vrmm?zNT@D&v+!N_7<{ z5Zw_;tZ1>N)KiyzcqM%(_DLIxR(w9$Gi~pAgEhC38MPWSF|o2g<#>=;eY~4AO3tyG z%-W1na?_|nc5Hfk2ZtV8EVS_>)H-~mxn^pNki*71XJlkV9?;@`qa!b7BP*+mLgWV` z$cZ`$ahQ_FOueGLkfNzb#5I=@--cM6gyPj@cYR`V>WysO*>P85_Gaa_2q80bNEFa- zM0gw?gI%JncNNq}EW)|K+>E3~Jlz*9BsP3|!B|194rg~#Yr0>)Sa^)qP0n#gaE8r+ zO|Wyc#>W&fW!yVHZdZ74HPr=Ka0hSfV-S7nP&lGZPJq}v>W;Q5ohJP!$Zsv@`pnBwyUig_+U#>M~=%ed^ z$4VAaBNAcF9iwD~+g0KgjyxZg{HRk!6IOAod3N{L>w>X?YS-Lr5x@56Veif-)0*dr zww&_2wzaSlHTmHNM&08~`QjS@JAx^+7R+$c$E69?S}o#}kPIv4%jCL6gY<3jA0Dok zJx1UEO9aFba;tu|)+;(Zh@r9j>hJ2iB~7uo{KT@;BYvFS|6d2mg)TYJ<^go2XjbUo z8^I8#cVR)Xukr;MCNl;g+WDE25d zkDd6aR`7?>u|7|au;h*7MbHh!_tI2hIZ=fxf5zp_V)*nd>o(`KuyL#3dC8fKmzHw; zPT-&cDrTkuxnaGev&;iuMm@+YwnS^8xxDpH4Kx)0VdD1N6rOh{O89kRp%j9G*sH5^ zI_m<1Z}`cVM!TCYSo3pI%75MEmL zrDl%+bvXd8sg6)7M;7YNFXx_FbTN^eNwGb`#yMTGkX~cVhNUaH<&S2B6#D2UONmXo zpfM0$<2znR(3{v1A&ci9Cgx0Q*iO}l)xc*I+VXe~mF2F#DVxP;Ep*!mO8efEUv-V+ z{aR#JTH}5;$d!x!7100sD?G;%rJ<3-vC<#HmyVz9Rq1Ang@h5uv zuq31 zr=b<@cClXiUV!8{U~rE`Oi{5=SkL#%TEs*}Z}0EBJ3FDg?r~{Itdy3Uguj_Ytr7mb zI<#T_J6qTMGu?@Loia^S;|A+DcmeLhOsI7XsH)fp3Shc-^BN5Hd8~DUmrz{Z%#~g8 z;HJ0~>TPanBk=-V#6aHTvOJKy&^Wech_z7X=p%msjTo=RG6XcOq+pp({bvC6JC~h2 zNxN``>@XSm?E42?dR$c|)6nd-u1%Sm8Xx})CRlNaWg1y=<2h`gAK!qTAAA!!^0pMU zEyhsIG`{gnyOfH5>V^JYLp^~CZdn|x=5t(_7U}t$W!nQ9Pqyv7a;)a@(a%B25ax7U z+(L3XAk^=R1xnz71z12JhaRQ`HMvkR--C~7x|hTZqMj$3!U557f$-88tc^}h?oW`e za>H)=Ang`NQ)cGc>3a7;T>5HEyYYrzXH@yPoORPBPV+%^(7w0xyviVoF))DYSP5K$ zNC1kFR*en#qJBG77zi%-jF&Cw_cK*($!wFHe-Aa=$3AT4exoNbAXgd6FPqjjoT!@C z7VTX{R6BOGBR4kittNV(Qo;LO1OZa<<=~K@^V+iQm?V|NnQjbls?4 z?iBS%r?9bb=zRFEZ&2&u*wOaOfHq*nmG(cd#Jp_3cP}%&o*voQa(6<~fs_xxGje`y zil1lEA`K3^Kqu?}vGF(}#towfUv%kDQVRQI_1Mrx0jXCBR*Ep{;7%Hs!91e9gPUBH0%x=z)-EVkJXHI7T zV=+rMWl$CoLi2@K;Pdf+l%k|eyf()}D{Tu-=Ww=A^*Qr8wlGbgf)M^KTeyPb_ngkU zKbe%Qe_9vpFomluZ+bgL zyS}NrmsV6b+SvR(7(+-B%k%+a3{+!`AJzyoZR4iYv&*%0xOKHz=ey2L89p~*i}lVY z!z4;&;pEcs<)))Qo`@W?TFjK2ri%u53qxt>F_vNSr|t+lA{YC2RhFnoFD!$AV!?Dn zX3O)TuHI6k$2g@>K#S5#2^K{}{X9BXUgAG0&3(Zn4JCoGs1^SSZhzndp}s^ttNTtQ#^_Ya%)n8rSxKBqA^{5CsJVz#h_8)0%hJXRb@W zxWFG^w$hR&T2kT|n?zQ($|%3O`Vxq*NH6whDuPln#wyHJjweJz@YOyW1lbC=1cvh+ zEqc_I>5CO*&v|)4PNPtD-5uIroRC($<&ECa5oA3%X$lqgjXsD}AU}Ays1G!-_vZzf@U*z<(V{t@u(kIK&tbx<5 zQ9wbvS#Ukz0ivOgO3x<(MWmH=v8Gsf0~C_DoZhjR{Fy4%7ZMcYz5M%_kFSZ>{W_7y z6}r+HjEYNtyPjYz&R_VuJvJ1xcalX6V8N??`TAWUx9-DLyi81Xtv<-B$N(Om*y$q7 z=6kq(a{LMn9X+SVW?=YyWGXe}ytud4w%2U1I5zyb~ z`@^q2LQ0LVG;FHQPfH1*t|yX5^TqzOc;=+NRnb{b9h}P_}h_q!Lr&&2x z7C~@U*V8)Fr0ggP-@)^r1~Z#0eZV4?)m_EMW)dRkg{6*=IrcF244#HB3$T%i;=rnrKwXrRA{iJf#f z?qQ&(2V~I_UL0)U=#_XR_YLY}1>PGOdRsUi%k2JE_S|s=Rwf*gVtX0&?Qu(~gLKe4 zdNe#@c|QQsSP;e3-}w!hq8;a=1Csu82Yb0B+wA>C>S1<4FF7TeiSo!Y*{OX+&93jr zDQ08}&YiXF8wEpb4jsR#5fBi-gj_j*sB#s?l&^LrOPiCvB6iwJ&x6*JhSyeC(`F?m zB^B#50j)-g-;WwpVSnjMBqGh}FL5G@}%D&7flHZb~H~K5%or_~qws+YF7sEPgkIh}8^=HGnac|~ZTJ1PX z?bdwx_7se?_6k)TwtFq-&GWw6Ts$4fnRj*A`yC#ZS-3mj8GM2|>%M-zN7JQ%DblbM z`D!+@P#>ZSfe4kZ=H9BzRNLzB6FO>gzamq>Vc6T3+F%TnRD;@tPw{g<()OVdsk62qs4eZ)v_8*YAqe zhi&NW)(dXHW`j5cmR)KH@p^;D<8lgIO2WYd2{1&Im*1INnQq$ckNRa_r`g%*Xk`Ei zPBV`1*3rRE$i48E5Rr#k_&eFWzq55b4R#xlW|fp)RyncLUcg2heliDm6b}YTC>0qv z%gZQCo%Xayl;=&5=w>lBO1f6dWZ<^+VhJs%#t~`1SJO4?53CuUi_T%v32jl**+r4d zjHnkX{U+omGh;8kc!gggDTqVH9*k!z5T*MqQZ}Wwgd)t??qy6utdJl%;4G^`d!wi> z_UD6wf^L?B7$l{nI@F@_--_TMG7CNAJ>?uR636fA$bG8yQ{U2u7HGqKX@QvlEuB9~ zp*>L0frErJK@o-Ich@Lto9^hDVAaNZejl421BFcs&gK9Gv1jOu)2VK} zs|uUAZ)5A6nD*US_gI$qf}g#Q?gLHl>mS8_LSi1s7H_P5V5Ab&4GFNLy_dCq*HJZB zpKMpOTV!h@m@qBUnJ?AHHpkO&+VYk(s=sNpOOc|Ut}I5X{C+?hJs>X;eC6zldu(j} zJ5^6D>b9?!Rs%(q#ZdLqfGUdEX?nE#C#^w5`4y0VG3YebI~^=cPfr8yoH~as%)Yc; z4{Gc9jwHpjAUne$PUZC zbNX?-v`5uW@%XUDqpwIgS{~V#c=imv0Ntpf`1kX6~$e>=mD;7KNI+>bXy77%S3R=;bSGsrmYq z7S~LwkRSV^oBdq1Eg2cv%a<=#{bLN}KL>ip)kK#rVL|6Z0z*QV`Qq>QBW{@cZP$sW7?C>B;D z%0(&x;t{@kYg1a^)VuRMDWyJ>(OIW~meVIouc}{>o@CXt4{9e0#=#5Z3>&G>Z|i&; zv*!VY4!yw66awkg6^PQCj#yeD@7NZ+60OBf5PbJVmiz=(zO}V=b-hzOJ2h1R7Wtcl zCfuh^0{Dn!^2(a+qFyJeVRqN+megF`HF`S7g0M{>9kH;q1ZxT{?nUe`20Z;Pw4T+T ze=TJ?g5_Ggl~&uIpCAqV7o^}Dug}b?JUlV?*iX!mCG?PdJedq$QHAvN{)sr|B;!J@JkywCQ90{QQ6*bSQ=I0c>O&*^j(@ z`3+w|el?_xHJ|Uz@`<2#_>=NI+m|$mZ64b>ZTwwNxzRP&UUD>-9+hsis4JPt!MLl2 zsu)Ba)!^K~pN+xrr1N&eCOMhA_-gfhugi}t%V<@J$u44ocFraDjDlRQNLAM?To{1v zMTtQL5q~^YD<)D~h*HBd#k`e6>{J~o^KoBIhQKHyY2h=-X~8)gOnacd1q%B*QI*8k%sRTgcfdJrRqC^nybf*&+T^Woxb%^2z!xz>SxIq z>b2;RyXM!!#|QFe4bHYnj)-5;-8nwU+wba#3=hca8eRy z?gdWgMBijv3zgg?Ew{+hLgq>~U^2P=JZY0;7FvneiXAxAOmP&7kGV1 zz4ES~UpS?avkH7=;t5S(h^+tpo}$y2*O%+8Z5`XlAtRrNlUPYoCbH6SS4%qmpreMUec9`|_wT5g-VKFM4Y5rq(f4Q5JpD9&Xa$;&{}%J$Bb{ z<-Cl4xI0~(z7s&d@bQ?>IrP$1*!29G!+P79370)$TTvecdTS5wj-J7j?_=J8S zb8t<%Ix@%^h07d^yB~}ukkQXSXlSS!Lc{63e7o7d0rGS#;~IxM-* z2y=iW++hE94Sg*p!TL>8P&t{7y@}yl_Y|iZCuf4IQZNSX==N z;ZZFmnRnGp9A$zzQi~C8Wau0a*MXI@e8CEqhL9z~zwc z&@e(h6VKhjMF;%F(wa^x*Mjx+xEl{pI0YP^{EuTg7arV%lvs-phKXBrz(nW;A-6hT(nJe6A0xi>wQrp7)D6>ur*xni>;Heu}h z#N3xjT5$f{2VOC;75@yUTS3qC{O0wD4HO~!>|%GD@QfZ_;|uLiU}qq z0!J0_#BsE$th9g}8opZ~dQ%==@_o%yCx-<6rtIM;59WGpe;WM!oqjeCztyM=$sS+q5i39WAKn1)nV6~(%6tjnW7>$!$vXZ+Z>)4kC67q{v)Pxg4*+`m=c%#@`l zbr6D*HiHZkkK#<(NI&{t+71ETc3=n-ddl#S*-!_w1pz;5{`o4kXa)TnL#ikqu-c*l zw?0lPpuL^1vG>c6BM!1FepDfY6LvVxv}6!m9=S7jv1quW)(0x7`S+CtaoOBzq?U#U zjG+(v;;Y$3w>_TdJ}qnn?pTc)gr~hoOP^lwkW}3N4M9BpBrUp|Y_Z4r35L=bu4B6q zG^GiLm8oRK*NhFYYkb9umj+cz7g652Iz5Q=)tlCAjz!OMYRhy5kcvojl21T|fXKR% zw_C%}E{{s9<3JJTPgQiOrE#Tl^|;Td0w;_0!fle%0w|)~ocnrvFP}l0X_rh8A++%} zENuPeIZ4SnEVH~qPi-H10o~w)fm1~HA5u<@`jmr-3~KB?CAwO|+3WotJAKje#Xs3x zmDYS@w2Wj*PPmk$Ey;c2**j}X9Q9(%oDz}cemlbGF256!lU#-LDGOv|=4Z?ce}(Tr zYD>y{Oehp^#Ck>nxWwdqyb4SgGBQYmCu;oDE2O)wh;rwNio*-u+34DIvlyR6GL2Zt z+rR8L7ksA?-hCRtzYycF$oHJH(Y&?E{eCYDF#-C+po?^l_w??zUThz1;YIDi?*LnMLqojpFLXUs--r z`$MVa0q|v<$$*F~;m*HWKczeZ}Y{;^@L+WHfJqtbPj;@H*c(f-S}k@UqvqdV9R1zXu$y%>7I zeHG5~N*ruFI7nJoMGGo6Y~%Z@G@$7s;e13^=8_D&7f49I zKIlCHyS;Os5yswZY~9_pIS(!eWUL{Qq9$}PLjnw^D&209-*9?ee?~*L4QyP@ z?2mTR8(F050pAU2ZeFf*IzQ`3=1gjKI#&C9>vel#pu^rsl2)!mZ?5^!qo_GJ*W+Ze zx_p>{AePiw^!W7;L>ccjS;0h+1XlBNa3qb$PU3b}&?J$}A_|IsVTzhiSx#4w%xf|5 z4Tj3{xaffeM@CA?==yYvHi%xegrB~6P2Cid1D7G-(2+7SMga930e4P1*9)&&(R`tr zp@nK&i$^w(9+8TLjFtE%bcD}*tQ$>LUVW3AKe27$msK5)cl#)I5g$mG*Bn)@#Z%t! zhC`(z?$?Z6i_Mw>%PR`iNo<#{BBGScp_{LTnpm9r;u*Gx?wCRYig=!1F|t7knso(H zG6(FGy0MQp!D9fX1b1-X6e3{OiQO-Z<+_1<>O{n|)oi`)Y-tdgokU@?A}%qacyX`1 zqAr;j!}PQ-4QQKy!U?{#q{j)S8ki@3 z9EKFlHA2_jum5iYJV~d*&dFlZ+i%cyT~JCho>?Dr~_16fD-jnPv1{ z(B~HpS8=He(;r=A(is^@)*C7wM#t?fob|T9;LVHr5;S(mGUq}ElMCY7@Bdt|;p05` zQ&FKPe&3?z{6SEZhKH4%`O^m!Qa|ztg?oq(8r|pQt;NN6xSaAHI3b?9~vP=0kb5Jcwf2vAd zKuD!Ms6ll9tCL<7DD3_@OYB(_wR&KGOPg#>l;?HmcB=T2ngZyv zYqH7=l38Q=IF{29b%sfPj{}@R({Q*r0r-Pt$e;d2 zL(#YWg@zF-#hKN$@wVuPKYya4S!b$C$AMOrlsZdZ`d*#CE<4K{ldI?rzb#3b&X~5(zC{;C zQ871-0LX=ElbcVnJh~$@l%SyNpvm6TFfnnefwSd3Wy^849Eqd@fDavaCd5MD8NQ5K z8*=%SDm+;)%2W*3#E)6EyNc3fno`7aly-$= zlDfSboSe4%hgq7nksSEKVawpOxMX-o_XUtGh9BIA5m&f<&I|64!4XC@Kymi37`y|A z>>6r$b*f&JCkd0M@84J3w^Z|s_GDFAx9|HPrXBcl?vD)_pFHpW>L@ysS2YwrkK9tA z{rQ<#ioZ46pgm*3-VfH4nD=mnDS67F`UM=&(ysqyIDG5IX+{GX!?J3&YE*p;s1Cr2 zmx}6kK8~5N%2Fd`BcP=hpJPO|`jnCUhk1BVt(-gWgC-_6>D7|ov!bF}S2M&C;oMxh z+uL{wiiEUui(}((pHEOa6A_UP9qc%}+sh+5jmrA_>z%>DL5p_zv{Jivc@Lw}-tMsa zLvK0?=PqOI;D1eb|D}vbq-%|zjWr*6NRU!|zBYj)Ch;Uml!12luDey~Z+R8i zyf$R+3f}E3vzS#kh*IZw2fSH-DT%Wj8TOW|2#ZfC_diWQYWL6OD6L#}z(7lVI;Uff z`rc)Gv#Mpdyp73QqvL2$hD3DAyrSoovchCS<_lX-PAIj$kPVC4hqX+*kPzhWWcbEnBFIpKg8DJ&uq{s=kqyQ=wnW}ChE)&4j8zpq32fSBH>yGP=S z^^lUsSbciT4075h88qbQbC#O;e2Bl5pTBwkhE{3}R z;BS0y#>2*Df!dh~HQqnmqy)$anm&_PXoT*~fPQ(_r?}<){lP9dJYiqzQJ1AXy6Flj zFx3JY+!6VZn_EkZZ+(#S@Iq>JmDN~CgajutWfmFu3pMIhnT%zE)v3pgz&YH$MBmTg z*q}(SwTm#8D9AgLhQ^%5SSA5zavfFJ(}JG>^dKfSwv5OHG{MI6Rb0P*tv~Liq4niG zYR4)NjWUq;kcsT;S)xblx*#ZcdezII_T^-+Dp>k3n%$4qWiOtT_GfV`GD@C})5A`j zF}000a=#kD#pdQHDy=Lf?^z=2UcQV#M9p)KueN=Cb@|xqm90$}_gA8T;Aqg!!Acj6 zW*@fbLeyp4p9?Qi)dNuEM{-e6^6?ED+$&F}i39;BC;$usyOlpuV$S7;Rt@9+y#^6k z(quam5`BFM&k65w8xM51zzNp#-Wkiz`u3d{5fc(alTp;1IBs4x044S-Rli8;V`c|HMz74i=m{5ni|&eGm}N$o?na+ zY?hZ4uYvnsGfKmuw=jLI#SH1dfF-BRU;SI9$5N*^2eOwpv%9lh!LPYJdBI(a9~)E^ zm_^M+)ZLu#?TqD)0&6+(dxKy*H4HNJaeW*DCZC%|&nvP)o^OWWVVt%+V>6%Neu+V( zR-BUV}=YLzU2)OR&bFr2wRi9!CG+m}C>FV0o)XUq! z3?^tjYI{mPOY8`6kov+Np{CK()2||fW`N<30dtq#o54Yz%xw0Z$%BI;Ey9yf1Xch3 zej@~gLPZ6!hokXk(ZOR*ROWrKQ!^C(hji`*dJKNWg6LGd32tjz5YB5w_^-GJG8!*6erA zsE>_+5PE%nmv3)Os_?-asJn-fT;J0dy{VhBI5FxO%RL)qY}BMjM@I);&SW4+0L{Xy zTt}m2;eA==+UjTdEiSbo{$aElowAGo+4MTo=TVR3J5*D<#m+%EPw08Do0;u9Irbq2 zXukk{VK_2Y^)(Meu}jeJpHPejo2Az4vmFZQ_|s6&+cxg+l(xXNK@4bj4={J1V|zYp zw(x&uMACh<6Ge4)Tga}4W9W`jqj>U(h|0f`x!oP-`|Z`Ig~x`l&Bi5E6oMFz}} z40E0yfA1+Fa=E)+d;F|YCZ+MW>8KHl-(f>`yaaivuq>Y?efksgXTxoN=oAzbO;_s) zz#!CeuDTRM*mt7lHqpq-ZL|zl{jk3fjnD1RET&eoxxZ=KqMlBAgZv7uN|?-kiN@7TuZM-NzjE zl2;Nhq)cV{9YX&+4``?3b~g{Rr}41M{9oc>+i2)hj-9=5>3?`HHnoX9cfe)!dWL-^ zTP5pX0%4-&|Ik*bX@4I1^HvT_Hs%r6Vmk+n{&<>Q#>I{e zyl{UCXphysa)#`Aq;X*yOhHNC;GzrQ~W z#&rK}#%3GaX||4?6RiMsz~%YH>0gxN(~QSW(aLvKscZvj9iAJgK zFAa^!b*!d-psfv-Xx6sxE5YL&nu%z`!onmTQOIj+14lU4b2(GfEnw^@l2yW%NE;ME zQ6|>1T;eX~^iijy!FcE$)7rFgis6zXKRHYTQbG<|q3+kG6e(^BL26M-IIPwod>D>` zGy1FMjUk#;`m=+(Tn_4FKJoqEM!5+rH%VW`?;Sbq)LqNtIQ?>amm=etjkX? zujwsvrkc-;hhVNs{{)ZV^^y9>;YxlnmXJVx_~3MZfZ0P`pz;>f%E|qFPjP(rOk9}0 zIh}2fx3&tZsHgyR5|e9dtr;Ir1W1hv6bTH6wdw#ZGzYDa`Cm2M$zM1e`PnnW5z~yx z@LlrY-X>}I`om$I=-cb$8cjL1wQacGg-JXriPx6j-b3Mg$Ms^ZYwUbJ9%y!WrK$=O zb3~5&gS?eELXS&gc6OB~!pD6rQ2e%Y(1f@b&jP%V6I$oMQZ-03^|)38<=C6t z;IbQ*c;PhwD@J|BBzi^mMnlCGxU;|O%c`3;F;1HZ85;f3czjxpe;1ukW;{}vw3@d2 zo8pQ}_rpK0Vi>=mmFbH&p23tY?=S)LXbbd;mM3Z~=u5CIJi|3AV|(_jo9>{;tAJ_b zUMUqzze>#OMNN_;2Z7`Z!#08RisI8RFIqE33w)kz+sPA^mqcVvb|af>sn@ks?-_I9 z2pgsuO|GU=)GGcAd2Dx94cCr8Ob*e42~KE;5@aTiRoFMH;B`Ly+XAjVjfEjV7ZN6j zT=|TU*R8vDz-=vBMSOiVVU@ZRI}jYj&H$SWOw;A%FCNsta?z}qADhIA07{=BYaC`k zzZ1+Do#$}_W!=)aPEGGfL6Ov_76{PNKtzd_eC8v5+iir*wOvVC8ZWXKuF2Q{DL z3P1_%!GNZC-q)xUYqyT4g3~Ki)ci5?WB$j@FcV<0`St5ZSdJRjQ|xhEqF2sTDehf+ zt$9oanT7Vl^i^3U;}FR64sX9OA#LZi=$=T5i=D}0>&3>^bFfi}h2Sqyk5TeY5c#S< zs?5IR5erZ62X{XR*c71SzKU!RT9!0;&f8-2#Oo=d+L$0q*~K~f`nIrR)u0j@A+Dhr&fYu;PQ&-`fS7!MBWxlANwB{6&JJ#p3nXmP z60dp3#*IHcC_6-WTx)E=pf&LG=Wi`5)_+$qUkn_91UAc=0<|*G8I9CXuo}-B2OUn? z<3w;A07S^7wKv}Da%irIAJ`2hOXh{G)Pm-KyYAGKYJ{LzFy7kcruax!$>JL*E*19g zmUyr^|HBv=(BLwbbxF!0wclXy%uzB#Y4daWiZ(eZ9qwP=)jjD(P(iJtR;L-~H?(Q2 zK`g6;#fA^zA2-><^&`cI5k!s<)F?L@h8arPN%hJC(|EAq_e8F&tONxHnxHAAnqu3I z)t4j*mLRFi%L@SBfsUK)QY&JFVAsz#LdinV+AzLpx*JlAT@lD{1Uwd#urLwX#E&hq zh|KRRY~Q@y6iv(Fb~j!(RKiq>x`z@4n8)sH-pq8r`$0AVL2RTpqwpy5=Z%0 z2Ni5LjVtaqYjw^c?6#5>5FC-sHidqZnj ze>CNSMGON**`k0VfSKz>hI5QBCf0?BiZeAeMGE)=KvobF1Dk3yCi)x}k3aj@VuddY z*B2AeOn_M)_`N>mF-xVV>G@}=REHt^seA+o5`4ArkN|IFC=)+yqQc>3VUL#6(x z_;R*aG6VFbj*gB_T1MvnxR0R+>(Qe}+}zyVSmJ8>RhYECFpX;$2qfNsDK1O|h(~&5 zqs7JP>1b$H4!pLG5EPOdPht~S%C(Ai$E)i1fK71D7EafwhFX( zXl-ph+nId0TnPnhGx#Wgr%N=X=Nh>fBx#`nYe0s`frc%do@yu`At*u5W4DUr`pDnF z$Ax>oQ33eIyt<8azDj*3dxXB7$`hKA)YoomwYqp}Ey#C;9t);tmN=L^331_6vC;`2 zMvAB&kUryr6_o#MxI0&wikiqt&BYY#v4AGJF_2)auYWS{v{(q1Ht;0iwz7x>bH2U` z2E&}KYz1L_YW8#G<>eQEaC&=n0um84<>cf5^54zeGx&uo9OIm*1R(o@>J@c8HLFNdu`)}$^G5kvEolW7~g`rYZIsArqQE1>V@d_HX~aqs+acQHCTIt`!Tn6$Lv#>d5|to$q4>h}g!V1!Fc zOMw@9QgHA9Sc-?J6th??dy_`UZ~XTdJT`{s#jH$`L_p{nNeS%?F!%-}*HSpst8ZYy z&@ks2TzQ}II1uw40`q)zWaxNn)nIftH+Rx(se~NX+|o;(ClD_)J-!Uus}gZ4A3|?w z;~`>az%6wJ>{~F80n^swY+U8$Y^Q4kP2inj)u~bqMcJd*GE=y#F_usms`wJNi~v1<|%KPP2YGY%ghlfW7e{v7>#jIHC{rhgM zO7=I`GZ4trVVTr+FfoS{=)InTvdn8n026h96(7?`yoA_?QOoD65EEt`lVNX^FDmZ& z{yY~L4Pdwm=+$-L+<@2t!)i&xcIIt@{0TLZMMF)okiuO3dl1CXNc1fQTl^C4bqzSK z)6CQNU+RABzN{?krkho?Fn)Op97CUvg0!`bWY>@;67Z~W8lOP0k*(52V~ zuU(TYop$|HhapD%y7wz)U|bHLsCy)Qg`}vZ8s0zT8hkZcX83Ti z=-C4DP>cu(Ui?_tM#=1tpO!p!eMkvWbAco|Nx)!{Q6DhTdxXHntqk&EmQpJeIz&}Z z=9}@4SwIRFunp$2TkQsq`(VCKY#)GQtkl$)|Hc?V_<8jC3PD;Lp1@~H)FD!F{DHy2 zg(~Y53^DK?&NR3>VFt1=VYaPfqP%!AtJzA|#zsyY@|%r~jiKQ*0z@$E;Ku6C2J2o$ z>dnSrUv^QRyMae*2aLhS#>VF6=1_5Q&{TnknZUH7ztZkuh}wA~Ur-JQ|xGtYu{MK{$ar zyx#hSdj3`_0uxncHAX}(M#Q&n*K{|ggBVk+3as>j$U2JeNk zKCH%p0E)odk>95Ds*Lbs$cbP&Rwx;zAX1tj&Z{srReDe49eiRb{u8%dS zJAIyzVRvJdThqOqeN@ogTR>z%rT{Vy!Ca7)lG3M(dj`B?9X92lbd&;FB~`MKS=7F; zx+heT0U5emvnUW+)T>gEDMIrDUZp^NyxK>iKKe(%ym;D^B}%+LgCi<{6#D!70bRdj z3p~ORFNDY7Cv9I|Imf?t_P&gH89NYm7^5ao{fb0m9Y`5~BRG7E$bVE(TKa~~;#YpY zIedY}s=lJKqRFZ1AXCM8l~uM;LQz%&T*C3L8vc9rkf{T@*qkWS$e`qHRZcQQJDUFY ztC;}z0xd@-F0Q)Yzr8*I2G+pgQk(0={@&lee*hjwuM$wA=JTlL6>k4C*U*0kpfmn& z#KmCIBeL!~sFPDwRh4m}4KiC&qduBc!V}~aH;DkxrybyR{_x^B9~5WstP4nULFwQ2 zl3R%|ST!_?M>wU`nAwiT0KT(_>q1XAa4=XDPy!FP1tDX{ z_~&kRH1&wbY!csoNGlys!VAz7l`UpZGlv& z{1d8HZRm46U+1)D$DwuwO728H&*X$}E&zW9sYntEInX<@VP0!WJOh{{G)q|TUNX{~iRqpo>n#*3_cH4T&u;0GrC9bw}bNT!PTyvWx? zPO-11G%cVIFBf5tKn!8?0=PA}Wt`+j*6x~Dd~f<8@Fi&NC1$Kwhs=^PKjCV2 zkTGeM;2h=%ZqHW{5fQNM126T0AQA)#7pAzM!j16rn8OUf(6f_W0zXa0&9Us8>C-eL z87dRy{+f$uc7`Wxuv>xQTmqqLr?w%%!NCm;JUuC-Z=FD2POZ#n*{R*CBR?aYxn?ex zbk0x~C#((vS@*gnBKDtn?VWQ4r}~xXt{TVPF_n3LbOsi>{De!TYCamoUZsy}+(C2( zNE}t8CySFyZ~Xjsd0Eg(qr&X$bW}F-QCv8w5?F$+e{UNiR}y8Fl#~GfZfOwFgHHv@ zWR#@m0pWjU!2)h4_2n>9mUlW!`t6dV>bH56`~fTEm=8U4G(o6O+yetYFhIPni1?(C z(pY~GnKc4F9fUR+qsPblqZ`W0&-!Vdv|u*)E5?P5inBD>u5^OLnteIaq=-|>pUUd? zpjEXB9$gV>FBA+0_{5~nC^Vw}DhjGXdS|NNd_M(ic=HI`45hb&klkOz$2UGi>iS4& z>!|RzRI((+#I6s8)&>O;+J_>j_Wkhv@}t!7iFPNoxC})-i(de|_CqI(@P>9`X)BUP zhL`tbg_TP%QP7M&2*4z_c_#$Z8 zZOwX9V}$quPap~a3BCcRte2M*b7Vd+s7JnY8Ww{;krS)7Vs#(Ns`4kKI3En&X)Q)Q z(Q8Q%JQaMa{GaE=2P1FUOFoIcRZVAua%lWbFxtKkP$`I6IWJI zuI^zA4QWz)vkW=CJ2e5KdPEPAGhJDtsySEXWgWrE-p3GrEk(G;qBPRmgQ?^&RUZX|(q%S)|)7 zT``99?&1XxYnqG7&;BhO{Hv{H#`8j!n9n_~0jhgnzd$#k0U33&t@~59T9?m0iqiW( z8UELAd`M-BnZ9A;Ej`?-8g0eO)_)x|q%i{?_gHg7WjpX`6EGiPJ6$l(yuiSGyc9{l@ZN zf@BnxZEYW59nWRZcHnebXLdT&5pFDe6RZ>EAW*nm3R7D(7^ue1E5t)pkaTGiGDacm z*3gKL$TlZ8|7R*5sj}%42Y0X`p!$8GoiFufOB$^S7}02QC*a;xD^{~_t$J13VHa5G zJ9Ok}czIaW$my&|&8;S}GS2dDq_dc%VAiHoxk}F1iZ0_`^S9%b#BOl~+HuA! z273RyhR?rzU&bO@)I8GT>6p^cJ@EY-3u^YWUFFUy+pTia6KQb`!cD88_y|djY1b!( z0ccJn;gZRX+xZQ?a~&yxeSDUZ8L9H?ZY$S|?@f!=*!xbL!q5mO*2HMdx=}>Ig3hbcotMX6M zkw9gk$QltPM0*#46Lj=Eb3qH^a>0nXseLGT8~yB}>qhZ#1HQgP%$}Z~a|z{th(Por zMD;CigS=o4iRvMVDxN{993in#gy7084`F>(Xe0FDO@1Sn-PZo&${v1QVV?xF2gb;66M=`on6s1 z-4_;;b!lPTwYO4u(uMyRJ&ByM4%o zq)b^=W;;VXcY(RoeK`5vjq_7s(oTjq5@nkYmG>L%JV!JoFwt*x@k=T}%i7?L?Ekgg zo11cM+5>V|(5>oC90zvYn^IrM3|U<4(M8Yo59sEOJ1!p@J7lA6F}G2S>{5Sz8UrUQ zjz@(&CK%)5S%#CL-gd`EV|!|IbPSBaDVl}iJ>r(v&+&R*1K2~;RnDzzicASx zJL`pgFvtkWb%t4*w=wIda0Gob#8PmwBI27B z+9`uLjT(a(2+j(ks&he~i&2vzd|R5A&{vrPV5ox)29SRP2`c6235(7mZ>o z)BeK(q!|gv0DBhh=ud9y-9tJn?-C*)Zr#39KM*Om`ZBJHvH491J`VH|f-0A;MGFk! zQ#E4lN6F9fgb*Oqz%G7u*81&<6RpCtsyGA|4awC2V(WrW+=*RXQX=v)X8jt;9=0sC z&WJK@6F*3U-;wB?IIfZd^IR=szwqpn76tiur;l(KGMo4Ajp+yGB(8@w8F6OgM;BbY{nJ&sHg&w72DZ)j;wTuyKP^oaCp)4VE`8_9?!ESQ!%-kq_=jA&xKUH zAdqPF3Q3q-&FWaLYZ4dVI)z#?e7`^QVG&*bCtEx7RYN)@CpP28fa8qgFKVVLL>&vw z{RN3LcA)zz)2r<|etuuXS%tXt-4iW%kTmn~v`}EsfFu`O%O5S_!1s?9ZoDQgMRNVCLTC4^+g%r8; za`tDmXD`4q&ka1CAS*B4w9Y4UT|8+pPPK5k5#5zhKT}R{rk6u^<(oLjH_H`Wf7<{Z^he{2RA2ii@#N#-G`A88Sv(Z))C6 z^?6i=-u_D{ft38f$|lykSQuApmD5|O4(?hYnH%1xw--!vX_!z#u1elNBChhiYH8nn zb$fXD81MdGPp0e)$ERfE;-SLqK)eQsTaS?a3ZgZls8dbpL8@R;{o&V>xE>~rC@ak=C+9atjsRcVAS6#6?=l~5iWcY`N0LBXltK8_x=V} zW%Di?CKm$a^2!sNE1@;$_QF0HBB*QwL?-cIlTiZq@A2)KRO-H-sYZ;;6ssXRh)ODb z#~bwoQDS+SG(3FGem!Ry{g-m4`w7ZFHXn{wzB63&;wW0M_I2>8ceRVJ4{cE3<=CpK zz6J<2sMgT0ka(juO}@E;!^97e{u#Tzh^9v&7j5Il_l`K{o~W#pKg*u0pF0zdOJo*%^0 z>d^P}F84AU^S=A=xAfDGk0|R>^s(FCCk|F)Ig{}#uWqqV zRBYX(i$#Tt)wqg|BTw@v=Z+ZmYl1}!IXqHQL@9OBjI1?nA8u0#xEi&Y0xbkuywuDq zcf~ohbN;4e?B5@;re9LvwS+n4fl#^Yx5b)EA!-iEE>|C4muUnm1@kD;wQlr&e3+3p z;n3lWtgY2oj(%Hyl0&95dF=7z4!(pyezVQ-RhX)oxDk7zKhLJ*u66F|V{;diPQDdg zB|gI8jxM=OiKIZ=n!C&-O9);5TfLwst5(-46mD&EQ@T^&I2G&&qN2fvH2Y-B<37%X zfSi3MFmy;aNDMK= zNO#8oL)?Skd;h;T|5x|9=hZx%=Q*6U&)#dT{aI_Tby#&fV(++$x$!H^Fa^#n)PZ~G zoOYFU)68A`_Yf%^fem%K5$qoc64puNL^^hFIjFXq3A>Dbyh+mcHWHs1@7c|U-Yn$@ z>C6djnqAtESEW=Zv00D)rN5#bQ0)MQL+&w0LvL%724N#=pMzmW^X9Bf+dc-{!ScP& zPR|M>P>{H!{hUw{vn*S!n*A=0g7Eyvn5 zYyp|0fWRugOE(uM=XA@0c!$}<1Me@oyeQhu&3{wz&J^;k{@kCyx^uuW3LcC16QXpoy@jYjn*hcY}eq`|P3izWK>QSaI%JWw zBylZAo z3EPU1&ZTlQo_?Aq{dkKL1i}LE6Xa6&h%p7E#-O=UjTBOUa|kru=>s#X!(os$srtz? zuRiaCm=k^LB%6qS=EZz`xa%_B$f0J}urRaoiB=r5JodLC5S{Y!u<#0ld9j@*1+ih- zDp{m?wwwax;mAsyp?IY<=#*ct-p5cxAi{-BMH`K#F$1cfEhq0>@vX)el(png6ec%q z8Lro=UA#3s*`Hjp?~|EH8U3KQj4mLlW=t=hJX3K*Hs%$HSkcw@Q8$1y(Czf==1n2y zWz)-$%k)85gogVkOYbic*x@CrdJ@IrJ*@qNDJdSi`p{OF?lx*O{@#}q#5>@vvUP*Z z)2MTx#dY-cbo@Aw?Wb^_H|yH9TBMgV#(|r;+-R_(#&txlDfHh=2_&g^e-R+Onc3YCSD>3n1{7$8z z^JCTea9KB}?QBTjF*DW+>7{yeO9j>K=It+CJ-fCoZpPzlE5xWq7Y;Q@*Pfwx*ed1d zP@l+q@RbB}h=!0LDi_@)!UTv?DbNqLl&>k5;mT$PGE!XoraCDNI(!Q%(4B(K{k4OY z*itNO849oJ<3hb8iSXnAjE(^mtt{Q$6>o-VG&K43jU3=Cp15Tlui#vbLkgBgWJp}A zc{<^lY;3QxqKSMTGq<+b!#I$wm1*vubrA1YzuEWISJyoU?I#svlQXmytgwO(L0R*0 z(1)sZBAe?W?PG-NwQM0-aIpcCRM)){$s~OD?x| zSKVEvj)lh2-kY#wN(Fb^s^gH$o4@akc=HMbrab; zE@dOu^ABbpXrf{Vc&u2Mwsy#q-rfiUx(j3C6Cp3Jaeu`Y6%KD1K4G#IX?>}mc%gbz zL9wAH#d^1E3{-NlV4T>ExXJgmStjO6fK3kqG6Q8hU|*Lf}2J#`%KpI0S7 zOj|aqB0N4Ua*^B(g)^6bvWG9E7j?T&6-Lh&YiZA`nnXMuYo87$upT7cAxJx|ZI4&+ z^z%#m%;%|7c4oyo9Hz|P$+DMvL@3U}$H%;lD-y@BBd0qw;a?;^T1)~Q@_K`8GhCbq zK&ZVS%Fox=IUQ2+lnkAdhPwb0;x2Oo10zq>O?{X-u&DTpS8Xf@>BoO%&K=tHu;W|T zI;TzxK5L_nHE*Z|NkP#Lw5_0MhEWv(Mx)hKdkFv zjz_m1hn?&QT{Iu;I1xo*=}@F*%fRW@^%UiuQlVTmo>t>aoI565bJwv*Aa`de>UK=a zMX(%?xOGy&oxf%ig;>JVjGVni(jaK^T;!5DKZ6}lqO6zC%$-)o-KMjAAI&lE9Pek- zY+EGGU0}7u`6hSYnPl9bE0Q6OcrqU7N2?8d`^&9iErlf!!U9L*ojJOgEz=1c5H1@Q ztfIqk);PjuLSVSQn~&8ceqmT79lbPRfXoJSkd|d_Oeliz&V`UNsJCB_^*F95A&=^A z^SQ+xlyob+ZNDmzv;3JmlOofqvuzCKlyLW_TS{npXZRK(sO^t_wCAgh zJO0vht$`c?oe6xE@iZ+TFQ6TaU7eW@c)w zjaJ@nG#_s=WgUJ_cBJ5P;-a1at4e9snu!2~1#AP7o%ad}KS8_|RF|8{Hf~_{38o*x zVU*N5>m}sQVVZBvGR%sJu;nz&oZJkKRw~5=__I4Ia|*)bvZzvS?x|N$@>S$Z>}Gr@ zn3=RyEge>9tRBfxHP^kj8}Terz!m886t>B%V%_v|uN|?nQfZdOqQD<-J*><8$NXo()co8at5${c_JztzEiMW1 zK1*(rwAS2`Qv7eSiiwq0@!H^#Ac;bOyCCb%SYbop;qTfa->KfY`J0 zv(!d4>)YQFigO|@UIE)z4`Cm6SBOuoFW=6^sM^TAUkZ8bTl>PG!P$`LpDT?aE9)RU zjqtmka@6C_%e=b$L0-*58q~RctGAI`9%Mk5JY8+6fW7V|03RYO-5n5oYC6w0bsbc{@oewt*T7G5!;=nAqkNEM1(9Gc9c{dk{W&|DT`w zy-*4*tI=Nc<>}N2_-uZLx$H)1P6R#A<@Idvj_;}zuMg6vr3M_5<8P9eC{45h1@+_| zl$zaasL{liW%*up9yd8{2%4#{>q zW$k=zY`SXc7@HQ6A;A%qO60Bgn$^rB8ESeAyE60^xHMhdZOk>a$Aa}=HtMTcdUR|> z5}GW7w+2hwWe(T7%BLG5dT1M-_A8ncO|@3%1+&cbRR_;P+K&8J3>7nKU9%7H7e{v;%%68klfQWo!x*QqwU9z=pH_H3FF1CJ_Zx*ltD1C9uzbj z&b6}Zk295)h>EHxbq4R7X?kPk)2pg^WczV!+3Z_fy+D4Rd{Vv49qe z4c?BaSqkzk54=__J-Q)OOc^$xS`XrD`Qg~xzA)Fix16T#P~syui@|743c(WoeirT2-gj+aro`ASF4g4Sat|4d!hbhAm%V8a99tOebgK?)MOwR8Mu+lGXcn-&c{AZ z1hrB}m7iU;&Omw}zvMaw6s{D+I3ne>ckqrL%Kv2`tnj8E%Y^6Y@;aLRTC0g^quV1e zjlAL*Ww+bu2&dg$eVlXI7m}$DE|5n8pRd%k*baTYO}E^9UuB9De0|+duEu4@l5|F%goWrN z0c^-GPj#--LltNN162 zSDNzKXtv2mz-c-c(qu|2Rv#Bu4MUYFuL671Xyo#h=}NI58d35>)mj3sRoO$Q2`NTI z8)QTG4ip%;33k9m^R6YNcMkSN$oE3#g-xV0bmFn#sQ7N$P1H`0o`QFIy_;}DywFHR zn)NziLa|OxrFP%6v*F=gk^x{T!GLZ|K@A-a%2CG!ij_yoB|S3BRE?+r)0q)R;PD%k z!O04rk@q82I$===ps(uQ$W=PrUfQeGr8F4(T?)dLbvm+1WxY5m;Syy2m<=TSOIuLu zIrjzdnJ49Mwp$e@TcaIrhsov|0`R}%$|^|}F=%riZ~Ob2akRuv>~-PQdbFxpe|CD4 z;yn<^VGe#+0-kNd=2x*(%Z*=P8m}Jc%nfo&SY088F>oi9Tissk{D6)&0y4zz_p;ta zp8@}t0oWI7QtrF!M8dc)kc#G#TMiyS(=;DNpge^jiEn#0;@%8s9}HdCV|IDkLXD!t zb;tQ*y_=~w51Z%Laf?b>``m4j<1pQWrHWAx_lfpvG_<<08U7~^u*_zL$|2p|e;@nK zBspHCP}QbaPibDuHs=@RJGM$u?z8tIvOzbPx*dR}hFIQ4U%%+LlsZ+E+UOO)97(%1 z{xNABiKu6rPPJvX%tw^17*7L6!{hiqe|Bw3bx05M%O>Su?Mc-7SWVHcGb-|()lU3! z3UL_88m}w$Jgz^IqE|`R8v?B7RY9%?VK+Bu>y{^g$QXY<%!{-bUqyBW(nc5BuD`fK zCw9<5F(ymAMyBuXV*9AaA?01XOaaD}D$AbbIUde7uysxSb#fZkF6|U!| zjCZ{Tk}rMlacZ9@1KTu=@9^RxY~?a+63BKO!SB?&Bpl!(4MneyqfCtQOY~B0f-~Ar z+t8QC59Yh^%GPhJ?vWDw=fME61Ey5e-BW-o=6`td|K+dmTS^2J3NujS$?peRfh# z?oW7?O8=$KfJoQ8oVQ1=PUU(pLjhR{_?$8hxh`{oo;?f8{;OB^Cp6$#a)`0x+In%B zrr^_yDTs@C&Gq&)F#bhX^o|s;m2GmS9&ExZp%4G&>j&OgVhUZ6V#qr6%iH>Qiv%KJ zjItU+j#(z)TFq@Po8Cv0-~D=rWF%%zPE&<=sq92VqLx5AF+Bi&^PGsbW(g7t-WJrO z6f!(o@(UoNfn;^hB83Pjf;|k5I;qndhqau6NB+0*=Zs&vIXvaoS>G0cg5oM{wnAZ& z+nL?Zf1NQt_10N?@6rp23OvUWhzecgH$2#OyWc$|Hd?$C^|3gYuckGLNmk>XZC74Z z{Frf$)jtsSJ&NFNyR07=;N+#)2Q0B$V~T`SnGo#5tOx=Mt)`=aWoNUQWS+{D6FSc- z%zs_Pt8C5`HN8h`RSE&v#=B_3d9P6_WmAcRoz$W|w}_g_v@2Az81Ile+GC;>S$J6j zoK?v7DG^3WUd&VmO8vI+v=qN|EM+}8C`&-RURr8lTwQn^vrWtQNt3EvpA;p z(!>#Cv-DO@PSN?bo^n!0#iChX=6}eNg*&kyoctdTqM3rI#f#v~s(TBa-gC z5FZ-Ib?Qd!b+M=VepA_kMD;%qq0@HL)pxt$Ljy$KC(rm_0zfc^s7wLb2{rw8HR(w_ zu<-zgv;`NK;$LH-G7J%8G!T@(wcF-9XuQ**i?h@d!G-?5gyl{V4T`R2F|xIu@sJFl z!OoZg(X*FyRE|o5wU4%!v*u;rLc8`UeymLCoyuSmxj2RJ>=du@R=;*rPqfmVIlfwg4%c+-5{F z-PV1EIfX4l{NLt5hEBCo3P?iCKos^@Vs7>pX(ec)x03>g?Lc%`Vzs{UYtdLoBH{*2x01` zhq7(+64sXJnZ0VHnu!2q!K~t}-tJBnPL{cPu!rZ82Hjcl2O@3!W1<>H32Ev&>`9W- z&$Lc-r8HLa4{iJESDt&${_!+Bs@J~>YBKj?4VzU?-P$8)U4cwRUw3gM*yLEebXaCH zlq`R&_-NPLnujcHPFww>Ef#=f292<~bc#fwB%zeE@XK=iRjgKurMQ(Emo-Gc_hX*B z0paq>NdC{8pWs@2H^xpYxBPr$P{nnly!41ls~)f;6i4nc?8)Y<(+a3R?#J92dzOU$mW{m$Fhbnw?;&gz3`k=x2j|M}&;WAMO9) zsQ*w+QVleo22J+fz1y-~*S2#Gy`h~yl(Me_Xb8HS z%2h@-no$Tsfe2wj^qc<$ZI?M%;JS0sUZ#09$HS|hv z0twZfS}6Miek@)_s5YY{Pb~kuQ-4MI{FSo~SVqPdpo#895`soD(@W>td(>F%fvtZH zBKTm{*YGY(Sv|MUOX5=zR`(By-Vv;!3vG-xce7`2O0U#z61=%2Po0Wp;gNubw-)&+ zBtIaGc3Lfmd`U%CS{zc8(Jp*^4coU~ti^=xRhC)-%OTAgoc4DLh;?K9A}T179BdJY zEy%)^PAD?}ir)+CVC*gUo7_sAeW5%P$Jc zqjAH66eu#tV!?Tr0_FtFKZ#q9TPv42g$x1y>iL4QBi}HO$J7(JT7Wk*<6>9{OQoX) zjwka>MbQ>Ww=q%3OKwILRtF@G?$1vH%qT4@`_Pij%6rtu;aSFU9c{vXw;JppmNaXU z{P9JzXrZEJi&XK&1ui{yethp$QF&9vy)HrdPa2f)$$~Qw*ZHgoa{)f!p#a6&^%7g$ z0vJjuY{89aVN`%Rx%v9(;#H7t@)62Kt}!&7Cx%?MfB$JSzZCj>KU>nEn&`9kar%S6 z%8prP#-j*Y^2Mi4#~|=~`C`iD8sFPfFG8GS+4bI9TX$YTNi$)K(T~Z#XBVi#m0lH> zl;}cAcGTvS)^=fOnx1yccliE}!sSJl@pbrOqr|I=s)nc5NVd}Xt=w5n^z=_D0q3_D zoLIW5l($YuYc2BUB!`tMTowv&kNiOo*T^*;$I9#7xonWeT4v>HPAU3=^)@n~y8SlK zkZ3eW8XdFbl&-o@_Zi`K;Nv@o=9pebD!4OzrX{0yUK7kwnwrkrFX+1$)$g`k(OC=r z;w{ZL`lhXwaauebJ65poCc!qXbS99bQF+f779m{ z!GXfJ2+&M4QPG}6UDLt&0{N$_XyUl*Qnf{|wtA#l%n-3~03-mX2)0~Mc^Pgbm+e|k zd=rryB^DKfM(|ioY~<1%&N^@GSE&SMaf^(y(jQ*761(x{$M8tG3-5OnMVwIztLBUf zRIN33e|6fOtto=zwJsTZ$=_4|QyUV**Q^U(k=NNEqi;Hbyz@Gx)^ec&%rAV@w|K4T3CSHV3QHvIobJiZ>ods zf^eFm)$k~?%j3H$*V(KJ8M3885#2}uh)+s!N|4+#+}e*%TrwWBYZ=;nwZCw2LhwRx zhVf+>DetzB;f|qbfk>g!rtbTp`)3atgryiv%@*oSj6PNuNn~I4Y`#tW$L@roJO9J( zs3;$YQfD&Ov*ZJs0rPCa8KbxeB;Z{5VVN1!;Vjo^^&maZzql1#FNPgZZ@cR3wD!x$ zeKsEE7Rr_i`JBi~F()?99rJQy$&uEa1{5!`TS&T3kiuXITZZUb+@4vUTHGs?(A>+cd??V|TS*v`7vc6NyR$<9L7Nn{(;i+^-nbV z`SHO%G9(_OssD%h4Xix$#|p`p|KHSaisk)lwb(hn0xu1Y4(G7QSs12!iQkf*)nr7; z#$;6$e#ZR=j3c}%35t|5G-oPd&cX7I71z7&j_yh$L@v--j4+~u(m=L3sC^7HfhtNF zE9Cb5OjY?%!NG}WTJyKfecsxy$5~7<^sXnPjuQDz2OF1~o|kT0L1_)I(3VA;wdLXe zIo794Pa~}KxZu>^w})eD1h7*k^4s#h(9DZ8ICW3>--Vb=WIj9Zs6(mQxsF~w!gkvSCyPgV6mxWn``~ai5cgN7d<+=+Rw?1 z6Yn2*N9UwRLoPTf1p}Ywx669yiscW zRs~ww0b4lty& z?|*`y{S`gq0BVw5+gw~p`R-QN%-_-O7LEGJG&Jrn%$Qb_jIkO0FL0RIE_1zNiEwh( zzkm(z{puy4IqnIfJ7H-gMEQ~PFA=l;tf8qgI`p{n7sgg?Y0A;enB~tcBPq8OqujyS zOG_hC9?16{sRHm{T3>)Bn9)-k?ebkC;QCDq)t+^ff~8j8>lN-5u=3@`UQPaGvV(n+ zY;o7w{bb?81CgqO?h@&&^mA-UEu^ZkL=ovr&8y;s3p8rkx z4A5V9Tf3S4g;mK)lveT?()udR9FZdDOH-{ru?Z|_x|XYHlq&w^1*?w8?X^lNOY>gZ zJORvy?gxTer#&B%6Nxsdv0f+z_&yG*9PP4BxGO{jGb@IyJasJElWhg#M>(i9exsG5 z!QVD5s5m}yZBY-ywac}fTld;iC~kea{6NTOk{mwU!RVa1-=zLP<0S-0D8R&~v7g8~ z6VucGVLh@Wp0U~V;qZ305&f5bSP<`TxTtRObohW@0PjifbeWvM+6HxxtQb8Z{m&Y| zrAp3C{Kz2TMWbh{JI4b-A4EBDDYQYhIe9qWSS^mn9@B&=QxS*sItl&>e;))1$^V4Q z89DueiRs>~)Pu^dT@X4euBn(`28x3m5z(9++XIbMDV^o7yLv@5B6Vcj!(xlV5701C z%R*CL>CSDW(agf1$$e$(6V$Yz^(%RjGoH&<50jFWZ8PN;F*Jy#n?Ss0r>|-) zN{=tz`OixjT%3qw-oHz&crqHfJ0vGCMl3Cv*=>%5Ur@cms^U`adwl&6`Sq+M$EL?f zd;v4T=Mf5`J%(7|KJN6U)*<;yBAVXU0-76KR*jY{_c(mnGB0Hg0{XkwjFy9|Iv7DT$F^J+n#y7ehaH62Z{Ff- zyc+&GbKUTx}51wm6>9i8;kGtN$2a91g4<>gnbTA z=_ZX8@p8?01NYBpmu8u*EFW&W%&6Urm`GK_EekkwiV?JD!dL!?HmBfO&BvBh`Iz{f zxG%=O9iBOz& z!$?ky5UFQ@+O;|KVz<7kRf?}$BC66{Axz~#PYpO!`AV4O4VTo{cOge({PBy^n0H!e66wP@;ZHsx{N8J^IsW7!kpw0sp)rTH{*l-V zI~nmR#W$+ALG~dViIcRieVA~s7b~QVG?l-|ICcF>*Oaxa6)0)NSJzD~f<`eVH`ibF zyf?=YWS(_A@SMfC^;IdbBuG*K`w82i_H?aY0;7Pj_fhQZa&?AK*Vgt?Sa^Ka&vl1K ziNE{=y?y&;A{D}@qqshO^z-Jvqu1Z3FQ;`9hgW1gzs6On>to?Kx2`5Uj$VAnW}?TO z^#vpG`SX92Th<);IXz02Pb-AKiCJ2^e|2G^bVl~5P^lvJnq=>t4*!S<2=3P1ta8c#``4>C@e!lhp@DKB1VCX;n zahJ~qT%82}9u55Z$@JeVVPHI0esY&*2wYTo|H%VD8CmN8ojCkGQ4~l##=!VSwz2Ed z*kZHyAT?H86+J}WEH(#b`0>r%Q^0<8(|;|ZUKbjmyB3y{7_zmS_2wwhg{pZ0G`NeM z=3Y=cV6gy(G5Ws_;jOK>`FTO!I55=QI~&>Tg^y}~{!m$ovOTB?doN=^=ko%0o)(rm z;}B^PrxvqE`w*SAS=^xp!LUb<5q}5>ftN#&f=;sU-KBJGjN{D_$=-=;iKy~-F z-)-u2O`M3?2X(H-gX1Zz)KiO#t(uK(jU^Trs>%fCKufv5%l}A;p8uVo?HzXXNk!9p#Vf-6iKdX(BZ;l_6He6K(;ZWDMt>&|z!bHf9y*dG}mUl*N z4FcG;V#&Ey^Y>Rnp6vIGKD1^M=o~AksS(e!xBGiSmVydiw$MV%&5W;}*{8$=7U!~B zyN_c&m=pPy`k#1TJZIdiTYPr$l~nn($bFqPYhO=xOA~hO?2je1nx%X)aOinJK08sO z*HjgPHi=jBt{Ykd3X~o1&Ew~n?pElfbloKRp?JyDpD=vXTp zmJZME1SLAf&MKVAy&{v9?<;x}XEX-sL%r11oIM5VP=4$}ar~QR&$E7r@K{+J@fExn z&zuUj<9^Y*S67!4I+v9!h9f?honE{>gcQ^l5wtn7L1j$01cX>JT8h<(QOH&VW3C0Ed9aybjK`p{Oh)^oax2cdZXy;t%g-bhknIKH%Q^Ly*O{o?}t zwU~kUNcMUe;&MGnTITifGJYo#soNL1mRr+}Op<4Pp!EKkYE&I{d$Zp2k|3ngA$My= zqn|~6+5A%$tIO&Jg@LAUTn{?U&vp&3K zDY4{6R4vTW$vA}BZ2Gj^{>7Dg?3<7EW?+MRGc{B2n;PdqT}!Kk>okK?|MbCf*@dak zYcJ#aleg0b&%y_8%FbUXJ>8DDkg6A&nEeJaF45|5PKFiqo9;d3mBLeYGKh<`O?=Ox zY)8!kRWFYCvJ%2GPIC()~3cEmwNn#U{(fBa6r;J@W^`9XFep{iwVE6Bj^C2EvB zd4gXnGpQ~={&)>$%iJ86K`4XG$)G$Y@>-`vwD4p;%izLsK&uZI9m-)5ywy4_?#nyZ#fKCg^i#|YIW|U+B{#}iqeEwVo=n)Z zND}K&@|5=VQ3T)o6_lO3Iv793p7h{5?tveCxuoX;$4MMEZL~`F8^nt~}UH zPb6vfTNOh3KeQV8>L_!uT*g|?EjYo)lNDJ<4ga%A-9NRZZ`xvanwA7EnKCv7D#(r` zcTx*4S)@o~iato( z&nKnDis+-%k27VEDdwN4W6ztV&0tD#oHTMp;}QE;HM&PwjmP{^ymAxu7-E-c+I$}){v;Ig}Njh@1x)^RZUy!9MeiGs!wLd zeZg1ycUu(;+JmlB?IlvYc(ffF_j}fP!b%y1!j5Uf4MD0v3E0Zj9 zjYlp+7@u4W%Wi}N@12$gC9j%UOm`)Y-G;sh`5lP(x*k>h0a=nFrLZqsgpqr z^_gF~kCSfACfG*UmtRn+Tj?adjkEA_Jv?z=Stl`C&Aipu>V8Qb#XY%Wrcy;2Dw2+l zgeah3EZ%SU^BT0&Yd$e@xodB2&CK>4Etu`u6NJ5xbNi=jFJ+^kY(8Z@5?~9&KbQQ{g^%8iCP_j{+>lX>dGj)3>v-#zpDo-bJ zRb8t*DZ~i5GhBa=jk+zbQ4xnSaRno!M({To$t=ycnqKp*VE^{?XMa)kF^=bI-JaW% z@<&V@8JN25d#!1;${!;h_t zrJxyW6=dH+7NTXXW)lYNDsqSTc@?41tB3FZj=o`^B=l3va-Y0bQ3KTR&&Tv3R_>Og z{a0BCd1FpJ>^@B{^WX_TA|+_n;dNZqbat?CU-yyXW(kUbQd@utV+nb~s;qH{U7f_tVuQ6JGPAAIysn ztoVc-gjoT)mOZ7!^v^fm<}M?wXtEnp)jfsstG+&qUR74MO+oD5X!$Y`-TF0m$EBd; z`y-yys`-7g9+5d&!p7lWv`mcme&Yir3%P`ZjPJ91#*lNXnL&XrKSV5J8E)yj^t(l) zhM;!3Dc;s?sf%s82D;>&e@^X(hh573Fy27Y-gjb$1WeD9K<@0bA;bu!;*P0uJwxBJ zz&w$o=QB^b4*tqHeKlN*cCA{WxJOucJgxe5d059OV3eDffZ0;302}@6#HmRvFzqUt z1+37gmzRk}u5<6uGsVo_M1et2A}Kjqfcbfthm60XQ|k*ZyJK*X@Itray3=N@f8w5L ziQkh}aT>kj&RG#Nt!3D8M%&}U2)=TCJ2d5kD$sFv@-2}wF_9> zKj4N?0k;&H9Ae*7%{M=jzRMsDA9Xn@kpvNt6g#QfD94fvetsTitm)K0HKPjgP_qe= z4x)empNI1eNQ_NqQljEo)h@)IIq$ZeA@R6=G6NbX%JX75#|mj$Y$JPVcySYG><4Wx znRPUFUTh!ia;5f5%YG)vk+H}phJO+Nv~;IUx>v8&fP9vP6kcY5WkAu)N~5V+gPiPS z$bMl9Qz0-VGNQjR)ebb2V7pk=P90ihIX)Iv*67-+8{T6Z?A2iUYog1;;qHt^z;Ei0M=NM{@-8(9&`q@ zeoG~jr#_Ow6g2MRlE5H5Qr#JJklM@Yrub#DeWe}0LRh<+PaA2XHec}WgU=gVQnD`r zraD^8&w~9lH|=MJ7@dq{T;I7qHv9f-5~oF@?AZ-5!J}N(lHDfvfp3<}>lq=M=0w(< z)7Eykq%ZsEe4y2r9V_gDX7*>AEHflcpu~GruryEKaQPU{4~P!9s&XB|X3~DPv+PsM zgMS;VpYOgAFn)d^_lmn(!$(W#K>ilJ=DXtyy$n>ymkBj~W+<(H!~rL`7mHrw6*X~Rs{XqxYu;db`L~|M6mPUQ{??Ga z_g5Wh`c!saH#rw0`p&g-R_ncO?#kkMzXT9s>Huy}KjvG??pZ(R9z-?yJ8(YVLrEVC z>{}e#ywpJVk={N?T~>VOei6Tw*Iu||5Ixp^t0nrJBqBdXk{;dG(VCDv=8)f7`VA-% z*MTspiaDsD-c*E-)YwL(vu)eScsZ&G%+T+KZ|{^#+f7w2KQ7F&6j7Ec<-~uekghPW zHb|{M`N-514Ym(OUdQ8|6L85i&&YK8-^Ypw95V z;3HvjrMTrie&QHL*4!eRf;~`XC2J!jm&=+3Af=jBH|_y}ZlT&WP?kTzPjpvHXCcVA zb5a$d)T@oDdP}WP86K&>x4@x)jOQ=Jf!a25=|Z~QLl6fETf5K3%@3Btc}Z&|O*LnV zE|ef5_r=~G*Kdgl&baEA#SdXC#rmI0HmTWvH9!9BnaL%d*qn!OEzmMA?51K;O0eo> zcq^tLQ|aa!SFju4CK%eVs_^9FghGP+&$qe%n5@M1*%aAedVl>gG=o!(P)N~&;J%;C zdimv07Z0z&ssuEd&J$^#?qQ2!YT0q%P1&wp3C3MD0EWh8un?urHd0LWAaS|CQI*AC zZ<=Y`B)^K=K-W>bcB1y!RTFcgyG*N$9Tb=27Ch0L$QW&R)L)m{EJ(}(E+U4#vS6wdu^eeTCu-cz7i&Yk|eG)R9>EtI|Fr$Bg_Rwc@<%|KaSH>=Ml^aoJT0^`X-?_YoVZ)!22 ztd`3cANNjKM+H)~UTU^@UU?dg$1t2mPU5R>T4qHb20rmYk&)D<&ExzqkIqM!ts0kK z(CI|1jK$D+s*$@=oOPf+`SM+F$YzWQtve)FgWT^=)!4jI5Dz;LBYL3BFEGuN17&E~ zt+qP^CP@PSWc<_Yo$BXS>=J;l_Y{oT+3PP+Q>%3UdTCgV&0SHiHh7MlwpOH$OZ^Hcl#rE z*Iv2?RHoz|UlPszoLJ>#ogk$@f1fCq*FhVXz(|TM21EZ)MZWKzyR&lc?cXoDGZVYj zEu=13H{S`k?Tb?yaG;7}6_D$m;5vRgo=b@;hg#@$s_i<**`vhYBnqqXd)%9}5fYR| zs;$WuUhaE4*3meXiM1-d6LEQy!WUJ|6nlgu#b*xXYcG>ZaFW)HnoF}<9h(Xu;$F>xlQtB2?j&T`8^3qt5OY8YKYY-dn zB(5a1iL4CI`~Of2C*upQBS@)}wl*aDd{jqu#&pZ!go1I1Bu)Ks4pTvDJK_vrDnkw$ z<#br`?+ot+@uk@3G}mz25lUy2pL{6KR!%L65hnI50@cHAdT_{v_ulPooHS?%Y?xcn zscSdTuFt3%mQGc!M zNubq%xB2mUXzXLGHaJ8F0@!%GHhXQWa9sB*LX8WuqAwIgueG8FSDq2)c;$ckYfrrmoa zss#3}1omkMzkxD~^>V=X#rTe%07#vyleoH$&d2QRTtCtV?<^WnrI6z6c z8vjw!%d!qEjHtlOyWAlRn-P1w85HXXJ55^pk;>!Q!%~!L`>XGVJ`s0e35>bp`>XGa z>;~tY&1CL%_(Y zanI~IpAMtLjl+7p?yKwf1jXpuWF+qMs zC+)+=-Wcq4NxtZ-fduFCU>2WT%(NUVeV+)kCfo94@4)QMSc9NF&EeI zi-`4F-NP+2W>7`4A~a{ngsIcJ-&5knxUu$X{yK_&3zC~iX0HB|k-=)&v-?p$cRhdS z_hTjDmA2V&l_SsL;^U(%dk{`Ze)C1Uc%z^$x#fzSR6(iB#TZ|17(1hOBTY;5)1_w| zPMnieb?F%r_-G-p1Y&K`a<@lRQ}cYn{q;-d?!_Mo|FsDW%6RaQ$L+vPkq@PH=f);= z7+D-86ZzU2MvGT@DMm{Ol_8VU%|`TTaR^nB$J4pbO)qqOp3bw24r7tWRmBb7kKDyY ze;FRbcgh{Bxu|V&?G~G^)ZNH4BIt?)ExxxS#X~ifVsMl+$}BO0?R1?zo6cx@#4Kq2glh3k)| zrrLBq+|niTwpB8qoyLVryn9t32mhUNYe--_&0YlJJg_|^b57xdkfg}UT3`p%cZ3P&m(Wpm_pd{)HIjq`5WAoL(TG@U2{Wq<=fxfHa3?J zo5TAx_pYajzeJOq)&{XDiHkB45V}G41%H%cWs)KiljnVIo^y_tt}Z#O&@L>Vr#-ey zI^5N{JY6}S66eNSR1X`8p7Rr~?rONVRVF8!Pt3QEd~TM@w-l%dxHC$FQiA^kbkLYK zGe3U6+o(!4`*tY5Jv8rrt;A|I)oyA-K67+#F(ud5^Z7LEu6v_l>OVw zg2IJ>4uB=pTZJW0%-Vf}ebA|19fl>I`!QAHt2A2DB~3Pkb<83Y0Y1ZrYGXV2>wM_N zb}@B_mO5yAOBIOdofi8t?XJp?>{{S{sGg<8VOEzYT+zekZ(F{9%7g%sgf=rrQ)YNT zi5zU~`8w|1uA9&oQK)zmqpTk(4RFC=d9d!B2Ijd&1mA3GjCK5WCUkriol9FIL7C3Im@ z63))tk~|F@2t6uvACGG9(%Al5psz*Sx2T8wfC=+l(!3vlOnZc;qzSm@dkKK+Ty~q(C$~udNrD*>d}N{3 zpVh{IJfb<*|8bKF6M}$TSYm!?`duv)Cd_0e4H;*R*3Ls1-p4Vt-PE~h(cKZk6!h%A zl|k8HHwbUK=7Z+J$k4Z)8eTG5egGvhAxW#<@u?TulR;MqtKvRTJR6_i|6OL}hW7K- zzLC#hwv`1t)B5}K!IU7M>-l1ZHsWNnR$@!3&Pd#rzH5V$^1emyd&p~$L&ah>wHWjQ zX6}gP(xuDQr-mtr+`%>pJS_~*nasztzH7Qlbkv#>$ab`0x zG+Ac-dU!O8kN_(Fr?In}Q5SVCAUi?z4@>dFnbJ zhd&r?9|O``jgJl;7V;v94=X23nO_jUTYLERT-~NUPM##9rg&P~IZnWg`tX-DH77mH zpLI`#+!stRuiPvyH4Rl*XIFZ}`TZ4u!vsuNQ?W$u`tj6$-SAQI`PlnIWMhuwZ^u_M zR7V!SyJ#m+zcs;x)nO~6iTF62WtkW@ElS*Fg|wH`+{TfnbVHwfX+WOxX8`l0-t;ggO|?QAaH z=0R>tlb1ZUFa+8jag~wQKKrd61?2N@lTz;g_$f(n}LI`Il)OEtG;x5;tw8cC5h@` zCiOT=o)@vYl~tUI*2)Ag?<|d632=h0v_n|Kx@?;s6O{2yaUmv`BL@V+Oz1!5@L;{g zIOuAS`Vg_T15c?2psdO5FEGpdSq5N1SOTNkxMYgaQ z!k`|eOVgA>Tp%^)<~PL&Xy3&b68naI=7cdIVD|+LPr8M91><_6Y~;)_kwwc$JWM0YF%31jZv-8;orYYsO1-uAe-<_Ds+=cz6g*;%j51Ha8){JZI~ zM7kP?sL|&r9R=Y_Fa46|;nxF7MAni_ks4&`4iG$F$v%tG=4+ip%O<`1joqr@SSz(( zW?}Ylw}KN|bbaDoD)T-g>n%fqi!i7>*Btlu>Lx0NWSnhJ^|sfZ4k=?J5H*e8jgli? znWkf0*#nYkEK}?!H8AqqLMXL+e=fsaHtE!ZI0S1>2rJp9KkdKiG9Zdo z8I=hHT^YYA=_f;{=F;U`ynV6oeF{7aVDt%N(~8ai z1TB4{vn8w(Nb+Yx04qb+G4~r$Sf+7B()F@xfFR6|^U1f!1 zul29Sq%K z@=r_{ELbX^!@@HU-E-I1KBnEa@_hfulcpZO7iU~Rhw%W7ic9;m&bsnkimK9IcL>*@ zve`EVbYQ-fybBABnr4XrCIKE}*hy{gnYPT)oAM!y+sb}A<8f!sBg4jjkSX=>=ApDt z<9ig}lFj(UIL%?{3`L|8jE^hII2+qv8)fFYiSuOMqUE-P^51NnJklGAW>-@$a7gE& zt4fjBD{J(q`!@e|1g9GfTwoFJLx=gtq7BwLdp|OQeK39V!$6`Dy&=R+!+AsGz&EPO zI>Vftz8w*%k9aDD9JN$4=e`t%s$0%C#c&zOp6yoV+x zO_mjLxnp_*v@5c{^EYDr)?ch;k|foU^q(Z(LVmL09WB%*y>T$YPtm_a0DILFuLNws zgQh5g&#==c?5xkQiTPF{_?CV*7}MWh-|80L*?l2^sl&dc9I#OPtVAhZg`L1;O$Y?_ zUa&Tl|9VCQN4b$9?+aTZKU$P7U;0J8Hsxp*nrHg%l{x3fnYZ6%E9vh|py;@F_JO;( z_|q9_?3X^RwS4!$0m*^xtXXW@Y2LsuIdklJQjm`w3136=x}kbZS{xaIIF#oeaaHiQ z&>>DV(`qa|IAV3I#9NUZ`14AN;@*Y8Um@)B5T-%4KX zizyryu0#1fg{cn@!r~Yd(q;_Iu98R&v(1Q*43HhPo;mLg&vXBIea*QabB{s5;4HeB z5HOQ&F8^Vg_xqQeHFmPZmB_81bcz+G_>d4kjuD{|p>U_dV-_e8BNG{a$d|YL`!Lz5 z)Mp^Z;m=qjhS;w`>)6FWB(r_HX?M@)Ouzjj#)e*Ob-lNr?Tn&-pI6X>;7FeY`7>~N za%j}~yg^t{b0zg}(t-U${5ipC*8j$`+uiadFR%yyx$?V@#zT^aMx!2f^n7yFhhq6x zwDFE;k=d5_6Mm}I>_Q&>MJ2F@hMYgW^O+|?Z26%TR~tHp5G`6be&Hu!B(!X^K!C z*p*|a866wSJ;lG4U7k%baj#$60nL#Yhuen&2DN~~u&|^KvBy2^DrnDU{pF$j?<}tF z%fkD((GHiWW~-B=qOGV4TqPs_HxCNj!op{POln?Gh#qnWCDMg5P%+(dDK1S|5H7F^ z_H*vO70v%b{rJ26nB=FH%~@QMQeVwGtp8j+H)z4VCgJt7dZqFt*^^a9Y z|6-z{YTRyAQK2MvTG<3^D3)+XabaKQL$GoEweftI&p&1F&4NbVpRgQ@CR!bue{6^~ zik`)~{X|E9O-hq$L`Rq0vUzGoBtPO@34hi%RqrX=)1;P!U`>btF<_%!5tVL&MBnQ+GxTI`S*`paJ2*k(AO&M8p-GWEm;-P0$`r(Oa)!_=(i zr4QJb^QB_1LBmTNK?fRfDl{^R7BFG0JDEob5R#b8Pn(#p7=cUz_>Dad7M?1N!KtYB za*{fbvEy~Qz(W4@b_lJ}=vL#Xr%AA(`SXN71kh+3%*Iq`J0FVpOAyL;ukG2@cPIh zMU1zq^zIM+Hjh7Zh^R}7TV2(xC(^=v{ksx_GnAWwx7)-}BFdT2sE?7q9=4J z)U?@;Y|n{4{|!mK4W_Q+1qZBL7Khd0c>|itPyblaBwJq1t}@+B^GYdsbz?94nq}Ev z`{y?-th%eOO8Ph~G1dI@uo`RvSX-R#s&11R!;lk?2CS+hUTeKuO~+EAXLk=gpk^#> zRUe;XKsvu0X!5P{2xoTudt1&*9GSIM)$hL- zqyMg~X?Vn(Y?p39!RJ_^qWeZV$9u=YiK{UL6E>Dc5Zb}t3I(zS)E@CQ=q|Z(Nr*h2 z;;O++I8Oc0vZI;gGJwJ)+l41VjGwN+Ry!O2OCXWl`7kL6)EOVjG9%rj)OdI#XlWs8 z_1O&d=Z@2F7TUE>WGlS|j;)c-mN|}sGF6jHi%X3pjIfWKS)Q$t^$AyhQ=8y=1vH9) zNk-8}z3GZKO2a>hBFBr~`}$mgC;kSwovvV)uT)8>Djv~< zk(1dgtmR-9?>BWhSN~$t@X<6t2~*ObMVXj zU_mZY8jOgY&S#kyPF!=7KV9X*>IU1rGSv~cfL`Wa(@L}@EN`PbFCig|`|-O$(sAh{ zu|W#LIfmL`gf_mFw-@DK5RM|{6&&Cl#e}!{+ZIa8U{EWvu;xeEw>8jvvev1raxKGr zW8N@gPrHUaC`p6vXIzfgZ>(6yvb+Q|R}#{jJF)D@-;(>N%uJWBzj#ONopK+ca^rc)KY&A|NI+Z_HZG2@XE}$ zW(Mra3_zs#OpGxt2UNF8?1NGz<}=yYn`I7<2`THP+mS@nT2zMe1x*c!#zQ3> zv`Q1bRWbz;)Ga7|6m)dOy!4n6m5tkl6iUN(Q&hQJJq~HdJYx{SKV0XND~7v_=!>*#4X1X2 zl()0ldFjb;V4gbBShs3fBV2E$%gXLpgTk9&@i--L8S=Y^Z``{Gu*KCncc9lQW6r93 z4+{!Bz=~8@=<_^@p1=Aev2@^^i>1VZ4u8)}6bO`;_bw6NNAxER0XIPOOR*JwA-?uF zOKA3Vt5|yBy+YN4-ER{s#|HiCd(i-l4WQKScrJn@yc-n51vg!Z=L7`h6*rYzv6b8Gr`6@9hDwx!=1esufHj2lg^tzEr@BqFBD!rt# z@$z=&TgagGWcFW6A0X9THn<$&G%H-(SM0BtyoDc?eK$mNciQ zr$=EKA1r3W@1mSnhWW?CZ>^DBk>#64kn&%lG(ORX_7$s^SxhT$34fSH8+%k%8N!6g z0{tI(UAZ=q3kk3}=`V}!UETxfEyEPzpxS9AR)_SPIh?Be8m?|*??2tzGx25|!=YtJ z2^jNVJ&wB(F+J$(d=12Gi&*oFQwYN>Yp$4KP>I1+A*{4{kNx9<%JhnM`_^*G+CaRS z;-gUtn%ad0L0P9a&5pT*c3qR*N{xr|vaLPJ`J<*1d6V#ynpd>G^c?i4iICo(C z>qS$08&u`B&igO=<>_4Lk z-;()TEL1(?!kNnT4xUlEdD#0g`Ymp^!yInzz=2*h0=D=T^={E#MdF7-aj@gopEt!< zK(hpeOSTtjM%8C(=8G6S*}B1i;{r8J`(gK#5VfU@ ztbd}Nh<2k{>t)TEsqcjDubBn^Duq^jG<)LBiEPuF7_PUo8zPLmrFq~QLVr3hR4z^SqDe$y6%isu}=yx?I0O@d~vVK zk=N-tcng_OI_*UA9wGY~tz7Ww7m+m~;fU9|2;be9Kat1#XBoKvP~_d0UYA+_Nu>GcCV&c9T}Crs&79#itZDeizrmj zXm8u+Vxpz;!wT9Js$VNCtQ?{D*oKT&7bepyJHyN=mt#4IU<5El*4GrKc}2+zS)Y1; zw11?u{CGM#ezLT+J-=>wc7K?L{YMb)^FW$%1=2=J^$GEC z$2tfi>iK-_`Ml(M?EZASJ=P*`{2683E?*@HOx{?2Knwf8zCvj#@_QivvUqum?kKFAzx8`2 z^%K^YHGjT`RYs3fJavHR*8t-Z(y0HETd<{aDFiVa~ zHp@z47Buvh)pL@I5ubs2cc>APsshQAdFH` z^rSR!pw%P_?L~xnN}Z zCUAdDnD9gUV##v6Q~qLF`62z*S4lQ`s{3BzQuzb{?ENNVeQ!e^1j*XpxQj=>L%^YK zs2SnDTQIbglH|W)VpY3-Q(P6=m+;X@xL9!0F>_s}5SP3RA{IDXMEcK~#APT;; zbEg0n)cPs#{A30W1Zd#{-_tsuQLZd{9_uT6+ob(Bzi(4u|0!v?LLP!hcIh8sL^kaN zt8)xS3BuMS0x@8TVoNN^zlU|D^yyA~0@tgav)Ufc>^HHTDwMsgIdUPzvKo1eX$e;& zsPM}O%ho#x!rYsjmkDC_y!B5c@Wj`Q0&d}>l}t-a6tKZKhSbCdiloi0BmX`K;hjGn zwLpw}`fGmzDEK{zzyC|GNHSDjqP~W4ugHIVQxo8u)QXZ?9(5l$DAm>{6FCPNR}g+# z+oc2+J{uZwmWiX)5c~h3Vs13))qjI}&ua{W$_Co~0v0nj`o={MXO2tP{j0dh6kw+ed;yo z-peD=uyDJcx@1D7;%Lmvsu}}Zvm9t*yhcL1tMM$&3D4I(u)UI2a=;L?mqQ=8m2-j9 zIi#Ak>bqF|B+sMmP{T3#k@N1lko`00okcsII|)9iinwwzQS1;A6g-9ah?@R+KKj41 z{?^{lb$*%=*3j;5JHku9{-7l%9(M1IGAyVknC|zsd}|vT5WO&wXfx?Q|2B{QaSIQM z-p?cP32f0+N1D>Iagtl~uNf(Ul_*zLy|Y@@t~jXJlvWzB3^Exq1=Se1HfWTXpv4!n zAVw}Nbq<493S6a&4iG3iPR(?>OH|i1EHeL!crX21+o|WzfgUh;GNmUm zlnK13WLzic$R}QDh3#we*j}~)Y!c|qig2uyQeSSB4U(6xeNEG?%l4mNYMH!sE>AYw_5&cobYR|)am1StMx&k^a@G-SbZ}_dlFj7_ zyXd)uvD5WKNZQC4R;Is1wT}Fp&(k4CsD-L}!vx|r zc?T#cEvLp878IX`r!@cKA0ye0Q&%_M}{8BSZ%*Y0Lv2C!_h*At#@e?eQpBM zp%?!l0I->4C660AZ1=;;XM4bA@>oGYdrRuGuAvZM{!*FHVHZp-a>4j!(KKqF2=fZr zd``fCmb)fhT77SeS2e|SIXj$-7NFj$z=6NKDoQB_J@S|^yL(y-iG$c97s5F6BP&?3 zRB~MhrjO89Vt%nXrIHQ9YRZ2xbdC-;bk@78Dd7!HQbY~KHy3-V@ds+KU9!B+559j* zGw_P2a2it|W)a{#JhrR*8m|n+Sfaf;*x~%uK4vpA?^l%sWEAML>%9YoVkj7X4RTg( zzl7V2Kd3Nz{;82u1o2?Y_a39V^LDqct_x+KyW#Kq7Qw0X_nw0S1&W$L`@?dW?PpU! zlo{WXBYqju;Su@GXX)T0!N5dvNp#F5b zLHOS(=dbe=lgzG-5H8!b3w2`&0v>0xs8~hwtCxWweLwTokkem53bmr2^VMMWOg5bi zW*0Y*!vs^qf>f5Qwe-Yxvj*tpd?nKhV!r3+wOgyU8)|Ys+ufuBv9{@n8z%|I#dsj7 z-Fr+uDX2=*@qt10%6%@6RaMS7p3OMl_NRSU;`x||jn+oplq6&A+jnpE=qJQFkrB;3xA`*c zGH&VZ}&S)ptNlW^@e5E>}Pd$~rIZ1yPnOyODSa8kUxku&bC8UXBp-!UYH4YyoN!XK-MQ+ihIIVTI3(gpkZ5UN zf-~M*$a_93VVyS>Y~|Jxrf?u+^_E~B_k-yP_m*h2;L*FR9R*l=H^W_8)-$cR| z{PziHztjqNT;5u0+Q5I6^f5S!bO2qzryxs6=pe==n70F!jLZh~WklCZCNNO^L(fy? zB2gMxFk5Z?@F!Z1yTIa_J(@a9c|>O#Z_q#4d9CSJt86#ie^or7mL(ViN+jD>Bc+2s z(Zc#vNF{2Ph@NTLYu&(USEMUVEq-S@?iZ<)r>AHO*30{c0lN1??P||lRexlLjEXHP(n`9nb9G7n$(-0i{d=0~jhM~u?zwmY{Fj)c z{tLB@%fk7GfHWM4l4UHD&^u-2mr`khZ$-AHq1tb1?Ma!}Vbn^jszH080$l8;f1H$> zGfBc?ffABqVO7V!MMC(#XSd%>Jk4!gW4w?bmnZzOgH+ ziXrbHvLYFl$#MJ&!=R?QSaYwc;SczZH`O_(Kwb$r_t~d6l<7WkCU1O6t_HuHpn?mM z=F8ue4|$r~wfxd=)%^yS9rI@Iq#v4ivWPA{NOL~uBJT@r61`BulI6?5kv!%x@_=U& zMhMH$1?Va2Jv8u4X3DxR+EiE=hM6xg@Ak661Kw}%O+{u@zd3UTo*Q|bb9MuZJ>$?= zR!&78I&Qt>3-*2q)KtRq(^|kCVhYtWk(%87X&eg{>2!mDSoo6nLb1%(XEeq^z zD^EoR?V(L+IMpoFY>&MA@XR)gQCH<0DiSsKYi|F6CyGtXjDAf1I@;Bsfj$DuD~T(2 zf$8x(`a$Zm`Q{ZNbI;r;=D$x11G1S}_eQ z2|BeZ!~>Z2>B`KZ^?MZ{Ud zZaC91QM8OJ5G8yYWjia^b9?>Rg*Wp-mGZxDN;l#{Dby0I8KryQ0T#a)7*iCUKiTXwYui6pu}G z%{FKyKCX7p8yZ;7?1KZ@eB-^ik2%00=Gj?(`fi)Uu|Ep8j{Y)2m?}d9(K=G;G?*7tv)WrSC@aB zjj`#9T#osro+E>pK;1#R^gr~%wRQEn=)?So=2PtyH%^kzaB zFdy7o4OUAN9*Zn?Di!3%5aZP!AjiKmue8uwvuI3`~zaiTG047X)+d-E)YfcdGoe`eaF?r=)H0qI`=zepePFUMo zqzq@q8s3s@9M$rY`_GG}NDugeO{GjNwIqVW^PL_}sMecjmU5 z+RtR&KG>=lHOP#AgR#%I$(>lOYDOO$?$T{?rSbx#wJ*6EnZGBt&Cby z#5PxbOB54a(lrZz`;UoN=2Me6S~(cP6~X2fG#`%5LFXNPiSMDq^0@&|RAtNE@~$2z zihiTJ!ed|2X15vuiyh;q@7#5eRI&)L2eOM9`zcpO3qJo{PGwIS#Wv|H@5(Lb99TM< zqP32GGf=Y|PotVb3~VR>jwhqt6#Vd!S2CVZiDr%g>GCa!I3n$0Oup*I%UJ7AcFJ9V z_L1U?iSs<&>`3SYmB@;gG-1C}o}c#&Km#hAM7LJv(?pERPf|c?TUnR<@NsE=sWD5~ zmg&e6kO&H}18T2Mbfd_6x_BrxDXQ*i8rdQejS3D?gOm14qzW$v9~_?d9R}aNUy!p< zaAa?TExR)haebUmHgF_36~StGx@-m`wmd)HB>f-@6RYs_+YC%>UcLkZPV)R5TcJ+5 z>!U_VQCHQO>>~_nlTAh(&r3;V*EfT+UOTIjwI`(a4d*xIIPvbZeJ}0G=$>5rZfQ|C zS_KtOmBsk}R~vYNEyqk0Qq#&~Suyd_SfX!~USI^hj!X+kCsFiO8+;2+92qbrtdU68 zHBv15s_k-zMJyhPT+`8_QcpWhsW6VbH27-61xjVp2yvXm@__n%fOexjN#3of;j{$&4-{3`)HbWeYSQ4oc0O+01nTeK zTpVQQNwb_%MNzy^IYJN0esKnFzF+EXDXEII-7mvWBij|R;)3q1?0g$oh_~xC{#Aqv zKNV}$+A9N)8f(dPTs|sFiqlAFCKW(6-o_bhZZBK`a~h7?VI8McnNx9! zb^hlmiD9BT-$~?}O75l49-tM47LJW65*k<#kdHa^@L25iO(Gx1H~ef7Q9uCySS*&_ z3;b=Qdzr1#O={OBpD4bugr?TqR_6A&|I8rWF2fr*A||>PA3%D3vz`esL*@>f^IL)@ z#mZap)OqjBcSv^%3-+nE%4g!#ISwXYphTX!-hb?@wI{cXOSk`QTe9<#``drhlT#Vy zBY4W))b@(1Mg-sv{<20+*M3{t z%{g}SVD4-}q0PPT1f$XSZe zipS+)s`=o1KwF6mn6lcxF0H$P_fw^{RCr4E?S(F+R-nQ_G#hfRRUW~L?U=R}R zX)t18p2VUc{tprzh}`?Qr9X4oNy|GJj)qZcH?TwqajBjP@@qiFf^MJVde(O{Cn&C{ z%2w-42T>ICodjpT=7b5$Gu$QHmWOH-8sdQSxiJ?MpNKs>?GXuC@~4j8^|7t?ElYqB z83#ZdG<*i&Gaz9Cn%@05Y~k(z&__WP>in`D*%Je29n9mI4CQ}vP|=pXwChG#f)XE? zr-U)Yai&%KRx{-Sr_0HpgW1(19qyjWl#<;K4jxp}d?$ib_+{(@Hu~rSDeZ@eCabl6 zI7m*nt3&H+ScdgbiVMxorn;JT>uDb)|N7f9RJ=GnqYS(C0Ug5qN?z2?)4;pPNb9d} zY`SMl2^y9vNKIrIxHF~f5n6NLKmZEoCdMEy&Fux`5htq(WlM}q>xV=m2!Lti6jcMK z<9(I1TZsd*MR0onEJh@SMm&wRx6m$JOIymu<}GlkwKG9--U`y>_%l1QA!NX5_up+q zf2htWg*IX|HK9PFt$)0W_SNx1Eoud()8e+@)Pf^ zAeQnPg&%-IVpmf|?}{a1MMa|h$_`+kVANEl41VRVc?rCYbD3D>tDnZX%~mjDZlCJ2d~h+r}<-Q z@^j(&vqkn2=k2r=x7o8sZ8s}p6Q-oN{D?|H>D|*cb&)uxNYT6!;5`^00$%`Mfi0c% z_1da|_)bWLNP@apJj+;&NVM9GIj0!DlP|EAAokAv+il9J?P=PTfH!b$^vR8L&Zoa7 zBwlc78ASKs&aw5gjV`mflDyJ2wH>ga?&X}>JLN_vsYh~vtJxu>G({OPuJLH zr+twz8H&VVK~Xc_sgHC4?9Qin=j7649`2M9icPhjfG6Jsx0|jBzL@wO_!w zr<^MA>FnvBNQE^2f#Q8G24-t7U;n3gyZ>vzz5o9~fBFCa4NCGavs>ku@A6!?CIT@B z@VVhnbxe|az zCyz|N5}%`w195x7I%w){6dty1pSBDlK>07~?VCP(#e(4e0!q_%4dz38BhkQHrYunY z0+SugU0(VvHH1kT!2`u5zRVqK0M^A^rX<5FlQ~x!9T!q7Q#=OAZ5R`0xOxFlL_&7{ z!ce{7@KI`h|1RmgOnwj5^O2y$YyHII8sEAlx477e>@T)MO}P*eEMNKAHk;?e>gYde zhwPYks-_M3{JW;Q2KGI%o+w{FT<9f$%82&^Pv_uz*!P)r-46y9glk=aw_0~}Vsh|c z@So3Sxg3-f=`7-ffi7EJ3Kb=)0MyJDYUb*2s96E^b70q0cae!j|c2NO*!?_wSoPCz965=<|fup7sU;CpfCfu>Ly26(CVs*8NIvr;aPl;{znu4 zRI&Q(UYDwhRj*Ho9YKzzOC0R#x$P|ox$5|Q)u2CuCk~hNy&N)NVAP#-&n12w<`fJT zu^Faah18x`TbuSNkPHANN}SiG9-}Rbi1?9;6Z z5TeYdLx=o8Hf<`)?ABnTPx*E|Zr?a|GAz_&8s|`8e;S4kn}?Ncnx5&yuOftB4=~LP zUnlsiMxtN#i4yJaU!XXgYCaK&E)OuM zx05^T>5AZlfx|dv-O|ZVM(FJS(Qp_jk=nFTJ+sWji`MtJm@+#``y=p&CBmMQ%S3Z% z@Zdl`1I4Y)pb+&u_P7GosiCa8W=?7Gudz`6$~~xD!gojV>A8VCQRC zZMyJgui`F@QS0gh9=6JydzB6r@|P6)^3LI?Y#K1s>DXZ98_7aBM`?@%&HmHBFgWNi2>+FFZ+W#&4z$rC z_*O&$Ymu!f;HuZ%0ZQboaZ$SmmHhI=+!IrcX!w>=V;s#HgDcFodR~G*`t2dVRe30^_Ze8vNUQ_`gnP#pJ_EuSUdkkhd5~ zy*E29(8i)Eh>*HYwsq)78)8_Am>g?lqYznRnW_{$vZ~${854_NY%|}8q)_1~ z@e&xY)1%!#G}GGt_>l$No6vC?tb1K|!ZU}?fDTdhL$2!QODH^2A&EEddc1P}@Ylxv z-7(Ci+*K&kr>#L4g=wIk1Pe4%O_^PL&ELBXae@OCtory*>^1Tw?mm#nd@*Cuo9CZD ziZ^JQLx*gDXcIF=)JjN;QNpO->yq+h7DyVnnV12PuJ*}Ah)k8nnE^wM&e}{Z968VA zoKq$$YZu1fL_y?5VLym?f7KZ?skN62yV;3tAxU+QcW+&l(@-AtA%Q~ax!*@c>PxNBpPm&p;J8++HDt>J+WaK#%eFpXEN)22n#a#73~LQn@P^=;+aGo zSKDOj*uO6f0OJ9~ypBNtn}?TlP)>ix=kAb+_IqP>}D$ieJ2BDw<`%-rXBl|1lt9&gXl9Lc)Yxe z0P^w0`p?k71|o0d&^~TYhV&sBw*?u*!HTqm_|JWXXGI2YWF+Zp%;c{M72B4OO%UQv z-kd>+oW2KzVa{)>&c1UWX5J7{mhC&%R4$6a|4I?#OMl=0b(TJNUjjwH~-`D-sAjLpiQO*68ab> z^-FkGvmAKR(SrRetZ$gc0ojX^Ini$i_Dn4veIM!zM^;Zzz!%kMyY6r5iSwip$^-rT z=^BNPWG9{Kejq{lmnGjB8Fg1A$+i;ny)Qai8*&nkiUgCbreTEBG3GYpZ5|B&{Td^4s~u zLjy%(HR7)KJ%GxHt~>DH%nghYzv@r;_`iJ&!txdoWFr@YUm(PDuN%v$xcyjjZ0GB! zRtWp_jLO+U0%q&;+zC5$lJ9LI#(U%I4}3{%(?lxHcP$Z`%ayE6#yUySvHyX z>0ysl-d`V5R*#g~SZAUA1!F%OrB)i+0T8BdM-k{AOylOOY0LH_d8W6$U?~5YRwT*_ zS(_FcpHJ&{zmWe2Rmv1A2=9~GespbK4zyMN;^&d7qLbBPKT~;Uhge$#{aUW9@ycQJ zY&lk3^+s`S0N85$S~w`iHFeoM1Jk92wGDvD4H+mM$ue{R$Ruv`xEYxl!fU%ILloA}{QEPuh=E zeM7E_On%Z{HU+*YekFI}v&r@A1$*0rInB9?#e`k&%7_qbP~q<5-x`UvxY_;APni~S z1a5A?+Pr;C-JYBfj(V_Ex3X9}>(d7{M+T;R;6&)`5rT-iD<6p};o4=H&SJp#PAI@` zjXRg$B*;RCq|i2`kG(1-PbGf;D!3gHb$`8}0_Pz}=w?w0_F4@Z;zIRyJ+?YgN1ZLz z>n<(2)ZU|#F{AziGrjaZ4$h595J708k%qvUpV_E=aoe2#_b)9{$C)oJZGYniNhBN@ zq*^o`>@WD(OU_!SjhO(3Qeg4=A?1%vbY>@WUG#9M#!rQW5f_y;!@HKVDG3^^BU7^c zW&wvdfdiRuG7y^6$(ED&glF5D(nrEu4#C4xqF>&F$pW{w(nFvU7vZ}Fx~Xc zqkbz{hSjO5Rl#!BrP_eM-s%p4$0G2Q;yJ*xyd27%G}$Zi4(Y_gVb;kKYZLwm%*3;u(mi%R@ zZpei`r=kUa9vuBGyZGj|M}WNtxThur^99MMH$I_`{VT~?Ubf<4O2rdjgbP~VBV2d& zK|~#D8;YcN8pG&Z0jzNisSXcmjw^j9A7^vuTuvRM!DQ&_bjK!yw!3mc@|N|CDg+&S z5Ga6t|2&KHT!DYePr7Pus8K&HA#B)Mdq_S^&E;@GdHbseE-l;5y(AW^T-ws3-c{pvwQdg^I!^PSu#y-qE>f-=sTOV>gGqYZ0E$a)m zqsd0&G8&X;M3sr$J|kd18Qfa)P)U9-j1gMkm@jQGosCpAd5%bhj&HXKCi$894w5A? z2=Et*%BeAPxD<>w4~Y6!bkP3f?!y~3-_^JvcB=Zj_NBK9Op56yK7`%MtG*i*C-+Wi zV7&DC%H66}$|4dGbAj-~I&K%Wvzl4b&kv(3t>Tt52$EErOvG6^@o1ei4w+VHs-?g!0 zXf{W@lX_&|0rW}*`>1e0@e9uj1E^H$S}evo<;71q1wyZt6_*{d@w-3CUt+}+*31$9 zDNOTLR)&BjF>f%ln*|&Y`l2c;?XrDrx0M8zBbJ?-{Tj_a<#(zI?Z#zc>Za)cW)qdf zPf9EZ172Yz(KkOyY#|2iT0c^BVc59U^cJt^PzPIxp;?yiUo;R3!(f1?P558 zC_!L?Y5C3j&7z?_EM}F~63!Q@FUUF{$Z!OMyuP@WJSMtpF{epRRs)0x6v|bWl!$Dx z4?VvI?hCrl69O5YYH>zIr#l)-=?LBx z<@w9<_gX_+;`CjV1-Kg$AOk5&5r2+pD}Gc%#bjO0Si$O69!W#Kyqqql(!NczNrF}u zJ1Yzal;AHAJx-h(*1_Kp9NDs-Zq*~SYjIXNA9WD(Hh~EpEl~2J19;ZtA_d$D1%9n2 z6X*vLAyE>$ZTzQw^hR$0E|E@@o>v;4`iDGNx2D*{^X`%pt~tED0rTORS#`um$xa5A z(&y*uNYcq{YK0+Ig7753B;}|qtKGrQNlwW;ALXDBYgH2v#xDfx16pdD)w}}O@TxtF z%`pvN(qW&4&enJNbyY5dS+bwB5+Ww`p)l)%)Fj3`0s9%AN=h`z^vAL(rn9O7R3ELb zod05v@Gk^)jGQ`E2~U!wgzte%-Ljw+{R8UKq+Ksl@Okt4GJxXt*fZDMv!D%9j#vpV z>}5)RLk^paPze2&dPpqeJ?N>0$i)KL{0dc0YQ31Tm(fdpf^UWe>V8t&FUEr#4UL2M zlW^vChqUe}xDHm{t2LstY$*l2d_YAMjIya$b3Fcbu12vjy&)jqD*j+a16mlGQ;8u8 z?C~@&;AnTKR6|45UPaNstO|vFJbux)tW2h6^uR9@{3%S4k#&|WlEb>b%kN1M89Woq zJc?T&bEVD-vfct?Xt-6tUVeFu2zJ}L@xI2`xN!!cSX3^D<>eEp#B7|aNzRuzdkn90 zjS)3StVR$%$&F`q#Muc?S_Vi4`=6TmY^m!vI$TBMxh^&%7A6(tFD)z$dupE3J_hLy zEdRJ8SKg2pGCOtI1bW7!oSk@J^SrWYn;`pgqKb5)5g6xoH$#5zMc`~?Ok(3uTdVl| z{^_RE$IKsicY790KA>ga>gg%rndTbg3zMY7a(99i-@ z+|9G%m!Wrk#0|#&Y!a4dO@g%Jf;x{u<3<)l8D*P>sh{r^E=5IPb^sQ;RZNVv&$#Cd z9FJCw?*{>@kU~UF=PCQ^aX>lE(4iP`tdMlCdAgp;bzacHf4OstArE6Y>`%aHZ5NvY z*an;`e`TbRVNI|!)aMBpARXUa{KS6V?CUp(RxS@~xh9fH`f2XFqB25fS!N8HNKC~a zx#R8Pxtm8d81UcDlD?Effj$$GeQTy$D?H_k29z*!PSh?R>xYjn;AxNV?!6yaS2#X3 z3uaa44kJH}lS0BXEd@~<_Pw)3Tn$iPqQDBul zEAOhS{?5toxu1fR)DN?q4=jys!b&KJlo9Qs2LdA zzYA|rwqCje!YJKa#jXD8Bf3$D#OCT2+~;t3f)y|$aPsIf_^0bm%?(7k^?x01_~&KZ zy7f0D(QlQS7Jn91HKZ$ncn#e994#np(bqoJII?aA(1C5zl)0;fi(Krt@ssMcV^|~9 zi|chQFaA2`$^QUdIx$=1;C*qdY|dU+1+f7 z;kVxUbz7jtmz`LjLBu2Q>eS5tTA2(+R{_Y;Gt~@v0b~tP)eX~OrJb2hdOx~E>o@rbJu;y_wvCIkPRU1eSJibJI zQF%G&l#WWh?D^TA^5#%e^y}PJ@Jo>ZJ3={u-$Ssq;8`Nev-Vj&yG|24FvvU;jR;TS zU2F{BGKMl+*OPdu<)%KA8lWolR6E9x#gNYz4@1YYHtZIiSE;0I_Q0bk3Mu{{woomXuME9^Bv~{e$}!*zsVug!qd)lg{ua7FZIOb9nR&kH*P^o; z3#S2Y0SijD3DA5eg#5a!+}2F>MMqOn^iP2pBDo8clLK*-onE@ATO+NnM#8TV=J~tU zJRV_o%)kyHi5s#CwA(%zUux`@qk2^4`^{_%6}i;teh0+p0(*&nhM{tE*^#c$eDAJX zdv#j0Ip8-AXVZO`{sTPjEEZB4_u>;poiB*=DC52ZQCWSQSjWTW<86z8qXs9};VMkh&N zMtolCpgG?!if>@!OW>{1rbZwkU$l1UVG@^0BC|%4J)d4Qr4xo#QH_r>D>m@BV8<&^ ztVor3@mAGZ2iRE|IITWlnepec!gkU5It3O>EZegO3>(1Go3}#P&ZC$D&!YWOEi^I>$&e#3MXd+a65U0Y=d2 zq>$=OCczf9Yb8;V;Y0ul?PHfI?UT|tzUq^{>9bLEP=;15h}y_&nReQ|<)@yNV3Uui z07#44!2Yjwl?{W)U9z&`a6?HgTA2QDx%IzGAM=T#n@65>U`2_AK@Umf96v)I$nA80zQkedIdO%Qy$(GRlQe>Jz!TntQ<-bwkmJMLOq#{ zTgr|I3wa)vs}N40XV{py!Leu*quo<>2c!!jDPRY`J!ssov1VL zj7$;)R$J1Y@8BQ6f?j<8*n=r^jjvQGGNVV~*;h=$E#hnsjjVnB(FW(GtwG^e;eyjZ zdqY^T__q)0K=J#D2c9=DELtyeVkI%prpyuIG+5F3` z?hg|E-#tF>q(Nw~w%&i{6;mevakzE!#w+T^lwn;=R_GOW+3(E2=48gEx z@S;%1;ZRIMwYgoJQV8cuc6{)qP7=7Aj=hRZomD>`_mkx9h*40ToMKwQEZ^qCk46@T z)2uX5t2qPU_Q?L;UdDJesTOzTW=u}OXrl%y>{fE140$w#%e$FZT00eoQe$x2K0bYB zyRJ9X8!kFFk!5Vpm5mq`UFxh_fwYQi9GC~0sy6o>6p+ss{E~?K5U~5Ph?n*Qp({P> zZ0TPkUi8~Pjk1>2WSBY0es-J>nfnYidSx*;_v%jtXhQG<)m@3sZvv%(DyY)5Z#JO~ z<(Ca+xZyb<3TwqEk`hhtm6uK)b6vWnEbb16(NXr4z|?G*leMbm;w7QXb??G}QSt;1 zcH=!(C(#S>YvmW*20+T7CVRGj2!nL>A)-`hSnV4YXjZZ-K*}Y6IseYb&3uS0%oI?w z>?UR$^S3W(LT(xcj-w^yqGP|Q_Y0rb|2O~9L{yDC7W&58KJ&_{yhfbM1ZH-f97;C7!F z;G?52Cj{S{Ak8C%wLuJigyDvE&Ua})0qU?o_G_A8%C|cX%Xe19)Fgon7ZKnsJ0G14 zeKPtqtKY*k5|9NvG5qF10`jp3`wz!(PaPkCk6w}ok(Uwf#mIG=+1afn10#NtKsP|aCn_~jq%Zn4g-;R? z^YWvBPkl0$1!iV>YQ1YE)V!ZLQGw7Aia${8X>1g|ITe|^im}Cg9mQ!i!oPHF(3!6| z_40xB$A*&qn2#RMX&%KXc71>R;X4+H)9OW}U^?RHvIZ0JEua2h|8W~#CKK%` zqm3btlk6sz8EeHGVOPu{|76+3=&BH*q%jIKXcWm zM_4O=l?0Z(>feKpd`v+8q~3B-1d^I!?h27TzizHh2w^P&2W^H|e?g_VD!q{ez=AN- z04(x^@;p0qv$YljK5ny=ErMp$kKlO`Z4?7lR@IpsZ^%kkWm*I(1;Br>iP^4o&Ui!o zz3L@op>~KSb)2|i;4FO`{+(68a>x_&QBG<?rbEjnZj1op1aOk6Joh1RG8=yfV^Yw^mUcP4yVK@ZZhb2Mi8~S$aA52I)-s zF_%28(IIvO;@4(JbzRrP2)V7}k$iSd=}(CUXc8IGP*TsspYh*^kN{+ZXw2~7h5k!; z+81Wal4Ybi_<&P|J5zIr3f$=@z3rzhPjSWdD0#*5s~EQlnq9vM#Ztx#xBaRgj%8Y zP+I9&ASb4fMRs-&AF(|VyE9F~MaU9hleeiJrHwq+i6?$Voi~6ap{rVKrB!NTypV1W z7;xR`lu(WIYb{5Hj=yFVX}uRl1H)>2V3OLJ!g3bM*FMbC``{sZ<(T>1b2 diff --git a/docs/website/docs/developers/general/contributing.md b/docs/website/docs/developers/general/contributing.md index cea66c124b34..4346f76c5081 100644 --- a/docs/website/docs/developers/general/contributing.md +++ b/docs/website/docs/developers/general/contributing.md @@ -88,7 +88,8 @@ this behavior can be disabled with `skip-llvm-integrate-benchmark`. The `benchmark-extra` option allows specifying additional benchmark presets to run as part of benchmarking. It accepts a comma-separated list of benchmark presets. This combines with labels added to the PR (which are a more limited set of -options). See the [benchmark suites documentation](./benchmark_suites.md). +options). See the +[benchmark suites documentation](../performance/benchmark_suites.md). The `runner-env` option controls which runner environment to target for our self-hosted runners. We maintain a test environment to allow testing out new diff --git a/docs/website/docs/developers/general/testing_guide.md b/docs/website/docs/developers/general/testing_guide.md index 90a81692332a..df26f7363026 100644 --- a/docs/website/docs/developers/general/testing_guide.md +++ b/docs/website/docs/developers/general/testing_guide.md @@ -226,7 +226,7 @@ To run e2e model tests in [generated_e2e_model_tests.cmake](https://github.com/openxla/iree/tree/main/tests/e2e/stablehlo_models/generated_e2e_model_tests.cmake), because of their dependencies, `-DIREE_BUILD_E2E_TEST_ARTIFACTS=ON` needs to be set when configuring CMake. Also see -[IREE Benchmark Suite Prerequisites](/docs/developers/developing_iree/benchmark_suites.md#prerequisites) +[IREE Benchmark Suite Prerequisites](../performance/benchmark_suites.md#prerequisites) for required packages. ### Running a Test diff --git a/docs/website/docs/developers/performance/benchmark_suites.md b/docs/website/docs/developers/performance/benchmark_suites.md index 5d4bd9446b07..0255a3e743da 100644 --- a/docs/website/docs/developers/performance/benchmark_suites.md +++ b/docs/website/docs/developers/performance/benchmark_suites.md @@ -10,7 +10,7 @@ trigger the benchmark runs. The results will be compared with and post in the comments. Information about the definitions of the benchmark suites can be found in the -[IREE Benchmark Suites Configurations](/build_tools/python/benchmark_suites/iree/README.md). +[IREE Benchmark Suites Configurations](https://github.com/openxla/iree/blob/main/build_tools/python/benchmark_suites/iree/README.md). ## Running benchmark suites locally @@ -126,11 +126,11 @@ Note that: - `c2-standard-16` for x86_64 CPU benchmarks. - `a2-highgpu-1g` for NVIDIA GPU benchmarks. - All device names are defined under - [build_tools/python/e2e_test_framework/device_specs](/build_tools/python/e2e_test_framework/device_specs). + [build_tools/python/e2e_test_framework/device_specs](https://github.com/openxla/iree/tree/main/build_tools/python/e2e_test_framework/device_specs). - To run x86_64 benchmarks, right now `--cpu_uarch` needs to be provided and only `CascadeLake` is available currently. - To build traced benchmark tools, see - [Profiling with Tracy](/docs/developers/developing_iree/profiling_with_tracy.md). + [Profiling with Tracy](profiling_with_tracy.md). Filters can be used to select the benchmarks: @@ -174,7 +174,7 @@ benchmark suites as the tool collects information from its build log. ### Show execution / compilation benchmark results If you want to generate a comparison report locally, you can use -[diff_local_benchmarks.py](/build_tools/benchmarks/diff_local_benchmarks.py) +[diff_local_benchmarks.py](https://github.com/openxla/iree/blob/main/build_tools/benchmarks/diff_local_benchmarks.py) script to compare two result json files and generate the report. For example: ```sh @@ -245,7 +245,7 @@ build_tools/benchmarks/benchmark_helper.py dump-cmds \ --compilation_benchmark_config="${E2E_TEST_ARTIFACTS_DIR?}/comp_config.json" ``` -## Fetching benchmark Aartifacts from CI +## Fetching benchmark Artifacts from CI ### 1. Find the corresponding CI workflow run diff --git a/docs/website/docs/developers/performance/profiling.md b/docs/website/docs/developers/performance/profiling.md index 11e1ca475bef..acc2f2212c7a 100644 --- a/docs/website/docs/developers/performance/profiling.md +++ b/docs/website/docs/developers/performance/profiling.md @@ -16,7 +16,7 @@ IREE. Refer to [Profiling with Tracy](./profiling_with_tracy.md). interactions and Vulkan API usage details. However, information at a finer granularity, especially inside a particular shader dispatch, is missing. To supplement general purpose tools like Tracy, vendor-specific tools can be used. -Refer to [Profiling GPUs using Vulkan](./profiling_vulkan_gpu.md). +Refer to [Profiling GPUs using Vulkan](./profiling_gpu_vulkan.md). ## CPU cache and other CPU event profiling diff --git a/docs/website/mkdocs.yml b/docs/website/mkdocs.yml index dec6aab05cbc..877ee45a9f15 100644 --- a/docs/website/mkdocs.yml +++ b/docs/website/mkdocs.yml @@ -195,14 +195,11 @@ nav: - "developers/performance/profiling_gpu_vulkan.md" - "developers/performance/profiling_with_tracy.md" - "Design docs": - - "developers/design_docs/codegen_passes.md" - "developers/design_docs/cuda_backend.md" - - "developers/design_docs/dynamic_shapes.md" - - "developers/design_docs/execution_model.md" + - "developers/design_docs/design_roadmap.md" - "developers/design_docs/function_abi.md" - - "developers/design_docs/hal_driver_features.md" + - "developers/design_docs/invocation_execution_model.md" - "Other topics": - - "developers/design_roadmap.md" - "developers/iree_best_practices.md" - "developers/vulkan_environment_setup.md" - "Community": From 9d382ff32c5ca43f2ab17df7b154286011452519 Mon Sep 17 00:00:00 2001 From: Scott Todd Date: Thu, 2 Nov 2023 13:09:21 -0700 Subject: [PATCH 10/16] Lint fixes. --- docs/website/README.md | 2 +- .../developers/design_docs/cuda_backend.md | 1 + .../developers/design_docs/design_roadmap.md | 6 +- .../developers/design_docs/function_abi.md | 86 +++++++++---------- .../design_docs/invocation_execution_model.md | 1 + .../docs/developers/general/contributing.md | 11 +-- .../developers/general/release_management.md | 1 - .../docs/developers/iree_best_practices.md | 2 +- 8 files changed, 56 insertions(+), 54 deletions(-) diff --git a/docs/website/README.md b/docs/website/README.md index 587beb714934..bf13a8890526 100644 --- a/docs/website/README.md +++ b/docs/website/README.md @@ -46,7 +46,7 @@ For more details on how this is set up, see (though note that the website organization has changed since then). For documentation language and style, the guide at -https://developers.google.com/style offers good advice. + offers good advice. ### Building from source diff --git a/docs/website/docs/developers/design_docs/cuda_backend.md b/docs/website/docs/developers/design_docs/cuda_backend.md index a83ddd88f6a4..74b2f6993463 100644 --- a/docs/website/docs/developers/design_docs/cuda_backend.md +++ b/docs/website/docs/developers/design_docs/cuda_backend.md @@ -110,6 +110,7 @@ dialect to LLVM+NVVM dialect. ## Example Save the following mlir in /tmp/add.mlir + ```mlir func.func @add(%arg0: tensor<4xf32>, %arg1: tensor<4xf32>) -> tensor<4xf32> { %0 = tensor.empty() : tensor<4xf32> diff --git a/docs/website/docs/developers/design_docs/design_roadmap.md b/docs/website/docs/developers/design_docs/design_roadmap.md index 309e5f909a45..7b373544a686 100644 --- a/docs/website/docs/developers/design_docs/design_roadmap.md +++ b/docs/website/docs/developers/design_docs/design_roadmap.md @@ -452,9 +452,9 @@ or `%sz = std.muli %cst, %dyn_dim : index` bytes doesn't materially change how the allocations are performed. Since almost all usage involves simple write head bumps there is no need for ahead-of-time memory planning or large fixed allocations, and since no buffer within the ringbuffer can alias we can have -coarse (*read: low overhead*) guarantees about the availability of certain -regions of the ringbuffer (*"when this event is signaled all prior ringbuffer -writes have completed"*). +coarse (_read: low overhead_) guarantees about the availability of certain +regions of the ringbuffer (_"when this event is signaled all prior ringbuffer +writes have completed"_). Usually any planning we may want to perform can be done in IR via code motion. For example applying traditional algorithms used to reduce register pressure diff --git a/docs/website/docs/developers/design_docs/function_abi.md b/docs/website/docs/developers/design_docs/function_abi.md index 50c225d61a23..1148dfa8d30b 100644 --- a/docs/website/docs/developers/design_docs/function_abi.md +++ b/docs/website/docs/developers/design_docs/function_abi.md @@ -12,7 +12,7 @@ performed in as similar way as possible in various target languages. In general, this requires additional metadata on top of the raw characteristics of a function. Where possible, this is done by attaching attributes to a function. -- `iree.abi` : JSON encoded description of the function's calling convention. +- `iree.abi` : JSON encoded description of the function's calling convention. ## V1 ABI @@ -25,24 +25,24 @@ results are composed of the following types: ### Value Types: -- Byte aligned integer type (i8, i16, i32, i64) -- Floating point value (f16, f32, f64) +- Byte aligned integer type (i8, i16, i32, i64) +- Floating point value (f16, f32, f64) ### Reference Types: -- ND-Array buffers of Value Types: +- ND-Array buffers of Value Types: - - Simple: Packed, C-layout - - Strided: Arbitrary layout with strides (future) + - Simple: Packed, C-layout + - Strided: Arbitrary layout with strides (future) -- String (byte arrays) +- String (byte arrays) -- Opaque reference object +- Opaque reference object ### Sequence Types: -- Tuples: fixed length lists where each position has its own type bound -- Homogenous list: lists of arbitrary size where a single type bound applies +- Tuples: fixed length lists where each position has its own type bound +- Homogenous list: lists of arbitrary size where a single type bound applies to all elements The intent with these low level types is that calling conventions can be @@ -53,39 +53,39 @@ these types, possibly by way of additional reflection metadata. The above are all representable with native constructs in the VM: -- ValueType: +- ValueType: - - Runtime: + - Runtime: [`iree_vm_value`](https://github.com/openxla/iree/blob/main/iree/vm/value.h) - - Compile Time: primitive MLIR integer/floating point types + - Compile Time: primitive MLIR integer/floating point types -- Simple ND-Array Buffer: +- Simple ND-Array Buffer: - - Runtime: + - Runtime: [`iree_hal_buffer_view`](https://github.com/openxla/iree/blob/main/iree/hal/buffer_view.h) - - Compile Time: `tensor<>` + - Compile Time: `tensor<>` -- String: +- String: - - Runtime: + - Runtime: [`iree_vm_list`](https://github.com/openxla/iree/blob/main/iree/vm/list.h) containing `i8` - - Compile Time: `!util.list` + - Compile Time: `!util.list` -- Tuple: +- Tuple: - - Runtime: + - Runtime: [`iree_vm_list`](https://github.com/openxla/iree/blob/main/iree/vm/list.h) of variant - - Compile Time: `!util.list` - - Note that these are statically type erased at the boundary. + - Compile Time: `!util.list` + - Note that these are statically type erased at the boundary. -- TypedList (homogenous): +- TypedList (homogenous): - - Runtime: + - Runtime: [`iree_vm_list`](https://github.com/openxla/iree/blob/main/iree/vm/list.h) of `T` - - Compile Time: `!util.list` + - Compile Time: `!util.list` ### Extended Type Calling Conventions @@ -115,9 +115,9 @@ calling convention (i.e. `!util.list` of variant type). The order of the elements of the tuple are the natural order of the structure, where that is either: -- For a C-like system where order is determinate, it is the order of +- For a C-like system where order is determinate, it is the order of declaration. -- For a name-based system (i.e. bind to `dict`) where no order is defined, the +- For a name-based system (i.e. bind to `dict`) where no order is defined, the natural order will be the lexically sorted order of the keys. #### String @@ -152,42 +152,42 @@ reflection attribute with key `d`, containing a serialized JSON object. The JSON object contains: -- `a` (array): List of type records for each argument. -- `r` (array): List of type records for each argument. +- `a` (array): List of type records for each argument. +- `r` (array): List of type records for each argument. Type records are one of: -- A string naming a primitive type: +- A string naming a primitive type: - - `i[0-9]+`: Integer type with given bit width - - `f[0-9]+`: IEEE floating point type with given bit width - - `bf16`: BFloat16 + - `i[0-9]+`: Integer type with given bit width + - `f[0-9]+`: IEEE floating point type with given bit width + - `bf16`: BFloat16 -- JSON `null`: A null reference value +- JSON `null`: A null reference value -- `"unknown"`: An unknown/unmapped type +- `"unknown"`: An unknown/unmapped type -- An array, interpreted as a tuple describing a compound type. +- An array, interpreted as a tuple describing a compound type. ##### Compound type tuples A compound type tuple has a type identifier as its first element, followed with type specific fields: -- `["named", "key", {slot_type}]`: Associates a name with a slot. This is +- `["named", "key", {slot_type}]`: Associates a name with a slot. This is used with the root argument list to denote named arguments that can be passed positionally or by keyword. -- `["ndarray", {element_type}, {rank}, {dim...}]`: For unknown rank, the +- `["ndarray", {element_type}, {rank}, {dim...}]`: For unknown rank, the `rank` will be `null` and there will be no dims. Any unknown dim will be `null`. -- `["slist", {slot_type...}]`: An anonymous structured list of fixed arity and +- `["slist", {slot_type...}]`: An anonymous structured list of fixed arity and slot specific types. If there are gaps in the list, empty slots will have a `null` type. -- `["stuple", {slot_type...}]`: Same as `slist` but some languages +- `["stuple", {slot_type...}]`: Same as `slist` but some languages differentiate between sequences represented as lists and those represented as tuples (read-only lists). -- `["sdict", ["key", {slot_type}]...]`: An anonymous structure with named +- `["sdict", ["key", {slot_type}]...]`: An anonymous structure with named slots. Note that when passing these types, the keys are not passed to the function (only the slot values). -- `["py_homogeneous_list", {element_type}]`: A Python list of unknown size +- `["py_homogeneous_list", {element_type}]`: A Python list of unknown size with elements sharing a common type bound given by `element_type`. diff --git a/docs/website/docs/developers/design_docs/invocation_execution_model.md b/docs/website/docs/developers/design_docs/invocation_execution_model.md index 7174108a32f4..c5d31ef7ba01 100644 --- a/docs/website/docs/developers/design_docs/invocation_execution_model.md +++ b/docs/website/docs/developers/design_docs/invocation_execution_model.md @@ -175,6 +175,7 @@ signal_fence.wait() ``` To the user this would appear as: + ```mermaid sequenceDiagram User->>@some_func: invoke diff --git a/docs/website/docs/developers/general/contributing.md b/docs/website/docs/developers/general/contributing.md index 4346f76c5081..807a734ec0c9 100644 --- a/docs/website/docs/developers/general/contributing.md +++ b/docs/website/docs/developers/general/contributing.md @@ -48,11 +48,12 @@ for some platforms that are more trouble to configure ourselves (e.g. Mac). ### CI behavior manipulation The setup step of the CI determines which CI jobs to run. This is controlled by -the [configure_ci.py](https://github.com/openxla/iree/blob/main/build_tools/github_actions/configure_ci.py) script. It -will generally run a pre-determined set of jobs on presubmit with some jobs kept -as post-submit only. If changes are only to a certain set of excluded files that -we know don't affect CI (e.g. docs), then it will skip the jobs. You can -customize which jobs run using +the +[configure_ci.py](https://github.com/openxla/iree/blob/main/build_tools/github_actions/configure_ci.py) +script. It will generally run a pre-determined set of jobs on presubmit with +some jobs kept as post-submit only. If changes are only to a certain set of +excluded files that we know don't affect CI (e.g. docs), then it will skip the +jobs. You can customize which jobs run using [git trailers](https://git-scm.com/docs/git-interpret-trailers) in the PR description. The available options are diff --git a/docs/website/docs/developers/general/release_management.md b/docs/website/docs/developers/general/release_management.md index 5e2e2fa89c4c..7cda3fb07398 100644 --- a/docs/website/docs/developers/general/release_management.md +++ b/docs/website/docs/developers/general/release_management.md @@ -28,7 +28,6 @@ request that some feature make the cut. 1. (Authorized users only) Push to PyPI using [pypi_deploy.sh](https://github.com/openxla/iree/blob/main//build_tools/python_deploy/pypi_deploy.sh) - * For Googlers, the password is stored at 2. Open the release on GitHub. Rename the release from "candidate" to "stable", diff --git a/docs/website/docs/developers/iree_best_practices.md b/docs/website/docs/developers/iree_best_practices.md index 101143bfa527..342bfd82d333 100644 --- a/docs/website/docs/developers/iree_best_practices.md +++ b/docs/website/docs/developers/iree_best_practices.md @@ -52,7 +52,7 @@ TODO: use the most specific LLVM target triple you can? ### Tuning compilation heuristics IREE runs its own suite of benchmarks continuously using the definitions at -https://github.com/openxla/iree/tree/main/benchmarks. The flags set for these +. The flags set for these benchmarks represent the latest manually tuned values for workloads we track closely and referencing them may help with your own search for peak performance. You can use these flags in your own explorations, but note that as compiler From 68a3f86838b87f71e15b1fac180c6b6a31336de9 Mon Sep 17 00:00:00 2001 From: Scott Todd Date: Thu, 2 Nov 2023 13:19:15 -0700 Subject: [PATCH 11/16] Rename best practices, move developer tips. --- .../{guides => developers/general}/developer-tips.md | 7 ++++--- .../{iree_best_practices.md => usage_best_practices.md} | 2 +- docs/website/docs/guides/index.md | 4 ---- docs/website/mkdocs.yml | 9 +++++---- 4 files changed, 10 insertions(+), 12 deletions(-) rename docs/website/docs/{guides => developers/general}/developer-tips.md (99%) rename docs/website/docs/developers/{iree_best_practices.md => usage_best_practices.md} (99%) diff --git a/docs/website/docs/guides/developer-tips.md b/docs/website/docs/developers/general/developer-tips.md similarity index 99% rename from docs/website/docs/guides/developer-tips.md rename to docs/website/docs/developers/general/developer-tips.md index f6a29880505a..04b8a670ae4c 100644 --- a/docs/website/docs/guides/developer-tips.md +++ b/docs/website/docs/developers/general/developer-tips.md @@ -1,8 +1,9 @@ ---- + + -# IREE developer tips and tricks +# Developer tips and tricks The IREE compiler is built using [MLIR](https://mlir.llvm.org/), so it naturally supports the common diff --git a/docs/website/docs/developers/iree_best_practices.md b/docs/website/docs/developers/usage_best_practices.md similarity index 99% rename from docs/website/docs/developers/iree_best_practices.md rename to docs/website/docs/developers/usage_best_practices.md index 342bfd82d333..5bd33a6f7b14 100644 --- a/docs/website/docs/developers/iree_best_practices.md +++ b/docs/website/docs/developers/usage_best_practices.md @@ -1,4 +1,4 @@ -# IREE best practices +# Usage best practices This page contains a list of best practices for getting the most out of IREE, spanning model authoring, ahead-of-time compilation, and runtime use. Treat diff --git a/docs/website/docs/guides/index.md b/docs/website/docs/guides/index.md index 99d800f70551..fbef5977df00 100644 --- a/docs/website/docs/guides/index.md +++ b/docs/website/docs/guides/index.md @@ -33,7 +33,3 @@ Guides for specific configurations: for AMD-specific solutions * [:simple-apple: GPU - Metal](./deployment-configurations/gpu-metal.md) for running on Apple hardware - -## Other topics - -* [:material-lightbulb-on: Developer tips and tricks](./developer-tips.md) diff --git a/docs/website/mkdocs.yml b/docs/website/mkdocs.yml index 877ee45a9f15..dc89682d94f1 100644 --- a/docs/website/mkdocs.yml +++ b/docs/website/mkdocs.yml @@ -139,9 +139,6 @@ nav: - GPU - CUDA: "guides/deployment-configurations/gpu-cuda.md" - GPU - ROCm: "guides/deployment-configurations/gpu-rocm.md" - GPU - Metal: "guides/deployment-configurations/gpu-metal.md" - # TODO(scotttodd): move tips guide to the top level "Developers" section - - "Other topics": - - Developer tips and tricks: "guides/developer-tips.md" - "Reference": - "reference/index.md" - "API bindings": @@ -174,6 +171,7 @@ nav: - "General development topics": - "developers/general/contributing.md" - "developers/general/developer_overview.md" + - "developers/general/developer-tips.md" - "developers/general/release_management.md" - "developers/general/testing_guide.md" - "Building": @@ -200,7 +198,7 @@ nav: - "developers/design_docs/function_abi.md" - "developers/design_docs/invocation_execution_model.md" - "Other topics": - - "developers/iree_best_practices.md" + - "developers/usage_best_practices.md" - "developers/vulkan_environment_setup.md" - "Community": - "community/index.md" @@ -256,3 +254,6 @@ plugins: # Some blog post names/paths changed when setting up the blog plugin "community/blog/2021-07-19-tflite-tosa.md": "community/blog/posts/tflite-tosa.md" "community/blog/2021-10-13-mmt4d.md": "community/blog/posts/mmt4d.md" + + # "Developers" section was added + "guides/developer-tips.md": "developers/general/developer-tips.md" From f1beab7122c647c04f152dde4200367ae71d9d0e Mon Sep 17 00:00:00 2001 From: Scott Todd Date: Thu, 2 Nov 2023 13:25:06 -0700 Subject: [PATCH 12/16] Rename files to use dashes instead of underscores. --- ...bles.md => cmake-options-and-variables.md} | 0 ...ke_with_ccache.md => cmake-with-ccache.md} | 0 ...roid_with_lldb.md => android-with-lldb.md} | 0 ...essions.md => compile-time-regressions.md} | 0 ...egration_tests.md => integration-tests.md} | 0 .../cuda-backend.md} | 0 .../design-roadmap.md} | 0 .../function-abi.md} | 0 .../invocation-execution-model.md} | 0 ...s.png => contributing-ci-enabled-jobs.png} | Bin ...ci-extra.png => contributing-ci-extra.png} | Bin ...oper_overview.md => developer-overview.md} | 0 ...se_management.md => release-management.md} | 0 ...se_promotion.png => release-promotion.png} | Bin ...ease_renaming.png => release-renaming.png} | Bin .../{testing_guide.md => testing-guide.md} | 0 ...enchmark_suites.md => benchmark-suites.md} | 0 ..._cpu_events.md => profiling-cpu-events.md} | 0 ..._gpu_vulkan.md => profiling-gpu-vulkan.md} | 0 ..._with_tracy.md => profiling-with-tracy.md} | 0 ...t_practices.md => usage-best-practices.md} | 0 ...t_setup.md => vulkan-environment-setup.md} | 0 docs/website/mkdocs.yml | 36 +++++++++--------- 23 files changed, 18 insertions(+), 18 deletions(-) rename docs/website/docs/developers/building/{cmake_options_and_variables.md => cmake-options-and-variables.md} (100%) rename docs/website/docs/developers/building/{cmake_with_ccache.md => cmake-with-ccache.md} (100%) rename docs/website/docs/developers/debugging/{android_with_lldb.md => android-with-lldb.md} (100%) rename docs/website/docs/developers/debugging/{compile_time_regressions.md => compile-time-regressions.md} (100%) rename docs/website/docs/developers/debugging/{integration_tests.md => integration-tests.md} (100%) rename docs/website/docs/developers/{design_docs/cuda_backend.md => design-docs/cuda-backend.md} (100%) rename docs/website/docs/developers/{design_docs/design_roadmap.md => design-docs/design-roadmap.md} (100%) rename docs/website/docs/developers/{design_docs/function_abi.md => design-docs/function-abi.md} (100%) rename docs/website/docs/developers/{design_docs/invocation_execution_model.md => design-docs/invocation-execution-model.md} (100%) rename docs/website/docs/developers/general/{contributing_ci_enabled_jobs.png => contributing-ci-enabled-jobs.png} (100%) rename docs/website/docs/developers/general/{contributing_ci-extra.png => contributing-ci-extra.png} (100%) rename docs/website/docs/developers/general/{developer_overview.md => developer-overview.md} (100%) rename docs/website/docs/developers/general/{release_management.md => release-management.md} (100%) rename docs/website/docs/developers/general/{release_promotion.png => release-promotion.png} (100%) rename docs/website/docs/developers/general/{release_renaming.png => release-renaming.png} (100%) rename docs/website/docs/developers/general/{testing_guide.md => testing-guide.md} (100%) rename docs/website/docs/developers/performance/{benchmark_suites.md => benchmark-suites.md} (100%) rename docs/website/docs/developers/performance/{profiling_cpu_events.md => profiling-cpu-events.md} (100%) rename docs/website/docs/developers/performance/{profiling_gpu_vulkan.md => profiling-gpu-vulkan.md} (100%) rename docs/website/docs/developers/performance/{profiling_with_tracy.md => profiling-with-tracy.md} (100%) rename docs/website/docs/developers/{usage_best_practices.md => usage-best-practices.md} (100%) rename docs/website/docs/developers/{vulkan_environment_setup.md => vulkan-environment-setup.md} (100%) diff --git a/docs/website/docs/developers/building/cmake_options_and_variables.md b/docs/website/docs/developers/building/cmake-options-and-variables.md similarity index 100% rename from docs/website/docs/developers/building/cmake_options_and_variables.md rename to docs/website/docs/developers/building/cmake-options-and-variables.md diff --git a/docs/website/docs/developers/building/cmake_with_ccache.md b/docs/website/docs/developers/building/cmake-with-ccache.md similarity index 100% rename from docs/website/docs/developers/building/cmake_with_ccache.md rename to docs/website/docs/developers/building/cmake-with-ccache.md diff --git a/docs/website/docs/developers/debugging/android_with_lldb.md b/docs/website/docs/developers/debugging/android-with-lldb.md similarity index 100% rename from docs/website/docs/developers/debugging/android_with_lldb.md rename to docs/website/docs/developers/debugging/android-with-lldb.md diff --git a/docs/website/docs/developers/debugging/compile_time_regressions.md b/docs/website/docs/developers/debugging/compile-time-regressions.md similarity index 100% rename from docs/website/docs/developers/debugging/compile_time_regressions.md rename to docs/website/docs/developers/debugging/compile-time-regressions.md diff --git a/docs/website/docs/developers/debugging/integration_tests.md b/docs/website/docs/developers/debugging/integration-tests.md similarity index 100% rename from docs/website/docs/developers/debugging/integration_tests.md rename to docs/website/docs/developers/debugging/integration-tests.md diff --git a/docs/website/docs/developers/design_docs/cuda_backend.md b/docs/website/docs/developers/design-docs/cuda-backend.md similarity index 100% rename from docs/website/docs/developers/design_docs/cuda_backend.md rename to docs/website/docs/developers/design-docs/cuda-backend.md diff --git a/docs/website/docs/developers/design_docs/design_roadmap.md b/docs/website/docs/developers/design-docs/design-roadmap.md similarity index 100% rename from docs/website/docs/developers/design_docs/design_roadmap.md rename to docs/website/docs/developers/design-docs/design-roadmap.md diff --git a/docs/website/docs/developers/design_docs/function_abi.md b/docs/website/docs/developers/design-docs/function-abi.md similarity index 100% rename from docs/website/docs/developers/design_docs/function_abi.md rename to docs/website/docs/developers/design-docs/function-abi.md diff --git a/docs/website/docs/developers/design_docs/invocation_execution_model.md b/docs/website/docs/developers/design-docs/invocation-execution-model.md similarity index 100% rename from docs/website/docs/developers/design_docs/invocation_execution_model.md rename to docs/website/docs/developers/design-docs/invocation-execution-model.md diff --git a/docs/website/docs/developers/general/contributing_ci_enabled_jobs.png b/docs/website/docs/developers/general/contributing-ci-enabled-jobs.png similarity index 100% rename from docs/website/docs/developers/general/contributing_ci_enabled_jobs.png rename to docs/website/docs/developers/general/contributing-ci-enabled-jobs.png diff --git a/docs/website/docs/developers/general/contributing_ci-extra.png b/docs/website/docs/developers/general/contributing-ci-extra.png similarity index 100% rename from docs/website/docs/developers/general/contributing_ci-extra.png rename to docs/website/docs/developers/general/contributing-ci-extra.png diff --git a/docs/website/docs/developers/general/developer_overview.md b/docs/website/docs/developers/general/developer-overview.md similarity index 100% rename from docs/website/docs/developers/general/developer_overview.md rename to docs/website/docs/developers/general/developer-overview.md diff --git a/docs/website/docs/developers/general/release_management.md b/docs/website/docs/developers/general/release-management.md similarity index 100% rename from docs/website/docs/developers/general/release_management.md rename to docs/website/docs/developers/general/release-management.md diff --git a/docs/website/docs/developers/general/release_promotion.png b/docs/website/docs/developers/general/release-promotion.png similarity index 100% rename from docs/website/docs/developers/general/release_promotion.png rename to docs/website/docs/developers/general/release-promotion.png diff --git a/docs/website/docs/developers/general/release_renaming.png b/docs/website/docs/developers/general/release-renaming.png similarity index 100% rename from docs/website/docs/developers/general/release_renaming.png rename to docs/website/docs/developers/general/release-renaming.png diff --git a/docs/website/docs/developers/general/testing_guide.md b/docs/website/docs/developers/general/testing-guide.md similarity index 100% rename from docs/website/docs/developers/general/testing_guide.md rename to docs/website/docs/developers/general/testing-guide.md diff --git a/docs/website/docs/developers/performance/benchmark_suites.md b/docs/website/docs/developers/performance/benchmark-suites.md similarity index 100% rename from docs/website/docs/developers/performance/benchmark_suites.md rename to docs/website/docs/developers/performance/benchmark-suites.md diff --git a/docs/website/docs/developers/performance/profiling_cpu_events.md b/docs/website/docs/developers/performance/profiling-cpu-events.md similarity index 100% rename from docs/website/docs/developers/performance/profiling_cpu_events.md rename to docs/website/docs/developers/performance/profiling-cpu-events.md diff --git a/docs/website/docs/developers/performance/profiling_gpu_vulkan.md b/docs/website/docs/developers/performance/profiling-gpu-vulkan.md similarity index 100% rename from docs/website/docs/developers/performance/profiling_gpu_vulkan.md rename to docs/website/docs/developers/performance/profiling-gpu-vulkan.md diff --git a/docs/website/docs/developers/performance/profiling_with_tracy.md b/docs/website/docs/developers/performance/profiling-with-tracy.md similarity index 100% rename from docs/website/docs/developers/performance/profiling_with_tracy.md rename to docs/website/docs/developers/performance/profiling-with-tracy.md diff --git a/docs/website/docs/developers/usage_best_practices.md b/docs/website/docs/developers/usage-best-practices.md similarity index 100% rename from docs/website/docs/developers/usage_best_practices.md rename to docs/website/docs/developers/usage-best-practices.md diff --git a/docs/website/docs/developers/vulkan_environment_setup.md b/docs/website/docs/developers/vulkan-environment-setup.md similarity index 100% rename from docs/website/docs/developers/vulkan_environment_setup.md rename to docs/website/docs/developers/vulkan-environment-setup.md diff --git a/docs/website/mkdocs.yml b/docs/website/mkdocs.yml index dc89682d94f1..7a746a5bc9a6 100644 --- a/docs/website/mkdocs.yml +++ b/docs/website/mkdocs.yml @@ -170,36 +170,36 @@ nav: - "developers/index.md" - "General development topics": - "developers/general/contributing.md" - - "developers/general/developer_overview.md" + - "developers/general/developer-overview.md" - "developers/general/developer-tips.md" - - "developers/general/release_management.md" - - "developers/general/testing_guide.md" + - "developers/general/release-management.md" + - "developers/general/testing-guide.md" - "Building": - "developers/building/bazel.md" - "developers/building/emscripten.md" - - "developers/building/cmake_options_and_variables.md" - - "developers/building/cmake_with_ccache.md" + - "developers/building/cmake-options-and-variables.md" + - "developers/building/cmake-with-ccache.md" - "Debugging": - - "developers/debugging/android_with_lldb.md" - - "developers/debugging/compile_time_regressions.md" - - "developers/debugging/integration_tests.md" + - "developers/debugging/android-with-lldb.md" + - "developers/debugging/compile-time-regressions.md" + - "developers/debugging/integration-tests.md" - "developers/debugging/releases.md" - "developers/debugging/sanitizers.md" - "Performance": - "developers/performance/benchmarking.md" - - "developers/performance/benchmark_suites.md" + - "developers/performance/benchmark-suites.md" - "developers/performance/profiling.md" - - "developers/performance/profiling_cpu_events.md" - - "developers/performance/profiling_gpu_vulkan.md" - - "developers/performance/profiling_with_tracy.md" + - "developers/performance/profiling-cpu-events.md" + - "developers/performance/profiling-gpu-vulkan.md" + - "developers/performance/profiling-with-tracy.md" - "Design docs": - - "developers/design_docs/cuda_backend.md" - - "developers/design_docs/design_roadmap.md" - - "developers/design_docs/function_abi.md" - - "developers/design_docs/invocation_execution_model.md" + - "developers/design-docs/cuda-backend.md" + - "developers/design-docs/design-roadmap.md" + - "developers/design-docs/function-abi.md" + - "developers/design-docs/invocation-execution-model.md" - "Other topics": - - "developers/usage_best_practices.md" - - "developers/vulkan_environment_setup.md" + - "developers/usage-best-practices.md" + - "developers/vulkan-environment-setup.md" - "Community": - "community/index.md" - "Blog": From 74a9705345c7839bfdf093fc5ab539f740c5e2fb Mon Sep 17 00:00:00 2001 From: Scott Todd Date: Thu, 2 Nov 2023 13:40:14 -0700 Subject: [PATCH 13/16] Fixup links after renaming and moving pages. --- .../workflows/ARTIFACT_SUMMARY_TEMPLATE.md | 2 +- .../python/benchmark_suites/iree/README.md | 2 +- build_tools/scripts/get_e2e_artifacts.py | 184 ------------------ .../building-from-source/getting-started.md | 2 +- .../docs/community/blog/posts/cuda-backend.md | 5 +- .../debugging/compile-time-regressions.md | 4 +- .../developers/debugging/integration-tests.md | 2 +- .../docs/developers/general/contributing.md | 10 +- .../developers/general/developer-overview.md | 2 +- .../developers/general/release-management.md | 4 +- .../docs/developers/general/testing-guide.md | 4 +- .../performance/benchmark-suites.md | 2 +- .../performance/profiling-cpu-events.md | 4 +- .../performance/profiling-gpu-vulkan.md | 2 +- .../docs/developers/performance/profiling.md | 8 +- docs/website/overrides/404.html | 2 - experimental/web/generate_web_metrics.sh | 2 +- experimental/web/sample_dynamic/index.html | 2 +- experimental/web/sample_webgpu/index.html | 2 +- .../src/iree/schemas/bytecode_module_def.fbs | 2 +- runtime/src/iree/vm/module.h | 2 +- samples/custom_module/basic/README.md | 2 +- tests/e2e/linalg/BUILD.bazel | 1 - tests/e2e/stablehlo_ops/BUILD.bazel | 1 - tests/e2e/tosa_ops/BUILD.bazel | 1 - 25 files changed, 33 insertions(+), 221 deletions(-) delete mode 100755 build_tools/scripts/get_e2e_artifacts.py diff --git a/.github/workflows/ARTIFACT_SUMMARY_TEMPLATE.md b/.github/workflows/ARTIFACT_SUMMARY_TEMPLATE.md index 48aff82ad1a1..c8c8df4f3e0d 100644 --- a/.github/workflows/ARTIFACT_SUMMARY_TEMPLATE.md +++ b/.github/workflows/ARTIFACT_SUMMARY_TEMPLATE.md @@ -44,4 +44,4 @@ gcloud storage cp -r "${EXECUTION_BENCHMARK_RESULTS_GCS_ARTIFACT_DIR}/*" /tmp/ir To run benchmarks locally with the CI-built e2e test artifacts, see [IREE Benchmark Suites]( -https://github.com/${GITHUB_REPOSITORY}/blob/main/docs/developers/developing_iree/benchmark_suites.md#3-fetch-the-benchmark-artifacts). +https://iree.dev/developers/performance/benchmark-suites/#3-fetch-the-benchmark-artifacts). diff --git a/build_tools/python/benchmark_suites/iree/README.md b/build_tools/python/benchmark_suites/iree/README.md index b5841eb69c51..414939a792e5 100644 --- a/build_tools/python/benchmark_suites/iree/README.md +++ b/build_tools/python/benchmark_suites/iree/README.md @@ -2,7 +2,7 @@ This directory contains the Python scripts that define the benchmark configrations. To run the benchmark suites, see -[IREE Benchmark Suites](/docs/developers/developing_iree/benchmark_suites.md). +[IREE Benchmark Suites](https://iree.dev/developers/performance/benchmark-suites/). ## Updating Benchmarks diff --git a/build_tools/scripts/get_e2e_artifacts.py b/build_tools/scripts/get_e2e_artifacts.py deleted file mode 100755 index 88754389a7df..000000000000 --- a/build_tools/scripts/get_e2e_artifacts.py +++ /dev/null @@ -1,184 +0,0 @@ -#!/usr/bin/env python3 - -# Copyright 2020 The IREE Authors -# -# Licensed under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -"""Runs all E2E TensorFlow tests and extracts their benchmarking artifacts. - -Example usages: - # Run all test suites and collect their artifacts: - python3 ./build_tools/scripts/get_e2e_artifacts.py - - # Run the e2e_tests test suite and collect its artifacts: - python3 ./build_tools/scripts/get_e2e_artifacts.py --test_suites=e2e_tests -""" - -import fileinput -import os -import re -import subprocess -import tempfile -from typing import Dict, Set -import zipfile - -import utils - -from absl import app -from absl import flags - -SUITE_NAME_TO_TARGET = { - "e2e_tests": "//integrations/tensorflow/e2e:e2e_tests", - "mobile_bert_squad_tests": "//integrations/tensorflow/e2e:mobile_bert_squad_tests", - "layers_tests": "//integrations/tensorflow/e2e/keras/layers:layers_tests", - "layers_dynamic_batch_tests": "//integrations/tensorflow/e2e/keras/layers:layers_dynamic_batch_tests", - "layers_training_tests": "//integrations/tensorflow/e2e/keras/layers:layers_training_tests", - "keyword_spotting_tests": "//integrations/tensorflow/e2e/keras:keyword_spotting_tests", - "keyword_spotting_internal_streaming_tests": "//integrations/tensorflow/e2e/keras:keyword_spotting_internal_streaming_tests", - "imagenet_non_hermetic_tests": "//integrations/tensorflow/e2e/keras/applications:imagenet_non_hermetic_tests", - "slim_vision_tests": "//integrations/tensorflow/e2e/slim_vision_models:slim_vision_tests", -} -SUITES_HELP = [f"`{name}`" for name in SUITE_NAME_TO_TARGET] -SUITES_HELP = f'{", ".join(SUITES_HELP[:-1])} and {SUITES_HELP[-1]}' - -FLAGS = flags.FLAGS - -flags.DEFINE_bool( - "dry_run", - False, - "Run without extracting files. Useful for quickly checking for artifact " - "collisions.", -) -flags.DEFINE_string( - "artifacts_dir", - os.path.join(tempfile.gettempdir(), "iree", "modules"), - "Directory to transfer the benchmarking artifacts to. Defaults to " - "/tmp/iree/modules/", -) -flags.DEFINE_bool("run_test_suites", True, "Run any specified test suites.") -flags.DEFINE_list( - "test_suites", - list(SUITE_NAME_TO_TARGET.keys()), - f"Any combination of {SUITES_HELP}.", -) - -EXPECTED_COLLISIONS = ["/tf_ref/", "tf_input.mlir", "iree_input.mlir", "/saved_model/"] - - -def _target_to_testlogs_path(target: str) -> str: - """Convert target into the path where Bazel stores the artifacts we want.""" - return os.path.join("bazel-testlogs", target.replace("//", "").replace(":", os.sep)) - - -def _target_to_test_name(target: str, test_suite_path: str) -> str: - """Get test_name from `suite_name_test_name__tf__backend_name`.""" - return target.split("__")[0].replace(f"{test_suite_path}_", "") - - -def get_test_paths_and_names(test_suite_path: str): - """Get the paths Bazel stores test outputs in and the matching test names.""" - targets = utils.get_test_targets(test_suite_path) - test_paths = [_target_to_testlogs_path(target) for target in targets] - test_names = [_target_to_test_name(target, test_suite_path) for target in targets] - return test_paths, test_names - - -def check_collision( - filename: str, - test_name: str, - written_paths: Set[str], - paths_to_tests: Dict[str, str], -): - """Check that we aren't overwriting files unless we expect to.""" - # Note: We can't use a check that the files have identical contents because - # tf_input.mlir can have random numbers appended to its function names. - # See https://github.com/openxla/iree/issues/3375 - - expected_collision = any([name in filename for name in EXPECTED_COLLISIONS]) - if filename in written_paths and not expected_collision: - raise ValueError( - f"Collision found on {filename} between {test_name}.py " - f"and {paths_to_tests[filename]}.py" - ) - else: - written_paths.add(filename) - paths_to_tests[filename] = test_name - - -def update_path(archive_path: str): - """Update the --module flag with the new location of the compiled.vmfb""" - backend_path = archive_path.split("traces")[0] # 'ModuleName/backend_name'. - compiled_path = os.path.join(FLAGS.artifacts_dir, backend_path, "compiled.vmfb") - flagfile_path = os.path.join(FLAGS.artifacts_dir, archive_path) - for line in fileinput.input(files=[flagfile_path], inplace=True): - if line.strip().startswith("--module"): - print(f"--module={compiled_path}\n", end="") - else: - print(line, end="") - - -def extract_artifacts( - test_path: str, - test_name: str, - written_paths: Set[str], - paths_to_tests: Dict[str, str], -): - """Unzips all of the benchmarking artifacts for a given test and backend.""" - outputs = os.path.join(test_path, "test.outputs", "outputs.zip") - if FLAGS.dry_run and not os.path.exists(outputs): - # The artifacts may or may not be present on disk during a dry run. If they - # are then we want to collision check them, but if they aren't that's fine. - return - - archive = zipfile.ZipFile(outputs) - # Filter out directory names. - filenames = [name for name in archive.namelist() if name[-1] != os.sep] - - for filename in filenames: - # Check for collisions. - check_collision(filename, test_name, written_paths, paths_to_tests) - - # Extract and update flagfile path. - if not FLAGS.dry_run: - archive.extract(filename, FLAGS.artifacts_dir) - if filename.endswith("flagfile"): - update_path(filename) - - -def main(argv): - del argv # Unused. - - print( - "The bazel integrations build and tests are deprecated. This script " - "may be reworked in the future. For the time being refer to " - "https://github.com/openxla/iree/blob/main/docs/developers/developing_iree/e2e_benchmarking.md " - "for information on how to run TensorFlow benchmarks." - ) - exit(1) - - # Convert test suite shorthands to full test suite targets. - test_suites = [SUITE_NAME_TO_TARGET[suite] for suite in FLAGS.test_suites] - - if FLAGS.run_test_suites: - # Use bazel test to execute all of the test suites in parallel. - command = ["bazel", "test", *test_suites, "--color=yes"] - print(f'Running: `{" ".join(command)}`') - if not FLAGS.dry_run: - subprocess.run(command, check=True) - print() - - written_paths = set() - paths_to_tests = dict() - - for test_suite in test_suites: - # Extract all of the artifacts for this test suite. - test_paths, test_names = get_test_paths_and_names(test_suite) - for i, (test_path, test_name) in enumerate(zip(test_paths, test_names)): - print(f"\rTransfering {test_suite} {i + 1}/{len(test_paths)}", end="") - extract_artifacts(test_path, test_name, written_paths, paths_to_tests) - print("\n") - - -if __name__ == "__main__": - app.run(main) diff --git a/docs/website/docs/building-from-source/getting-started.md b/docs/website/docs/building-from-source/getting-started.md index 1203e4374549..8ed175d03a22 100644 --- a/docs/website/docs/building-from-source/getting-started.md +++ b/docs/website/docs/building-from-source/getting-started.md @@ -153,7 +153,7 @@ settings can improve compile and link times substantially. -DCMAKE_C_COMPILER_LAUNCHER=ccache -DCMAKE_CXX_COMPILER_LAUNCHER=ccache ``` - See also our [developer documentation for ccache](https://github.com/openxla/iree/blob/main/docs/developers/developing_iree/ccache.md). + See also our [developer documentation for ccache](../developers/building/cmake-with-ccache.md). ### :octicons-gear-16: Optional components diff --git a/docs/website/docs/community/blog/posts/cuda-backend.md b/docs/website/docs/community/blog/posts/cuda-backend.md index 9c3c56c557b5..340c1d7eed2f 100644 --- a/docs/website/docs/community/blog/posts/cuda-backend.md +++ b/docs/website/docs/community/blog/posts/cuda-backend.md @@ -28,7 +28,8 @@ then shares some metrics and next steps. ### HAL support -IREE has a [HAL API](https://github.com/openxla/iree/blob/main/docs/developers/design_roadmap.md#hal-hardware-abstraction-layer-and-multi-architecture-executables) +IREE has a +[HAL API](https://iree.dev/developers/design-docs/design-roadmap/#hal-hardware-abstraction-layer-and-multi-architecture-executables) that abstract all the targets behind a common interface. The first step to supporting a CUDA target was to map the HAL API onto CUDA. We use the CUDA driver API to reduce dependencies and be closer to the hardware. The HAL API is @@ -96,7 +97,7 @@ can now successfully compile full models. ![Compilation diagram](./cuda-bring_up.png) The steps to reproduce running a simple op end to end through CUDA backend are -described [here](https://github.com/openxla/iree/blob/main/docs/developers/design_docs/cuda_backend.md#example). +described [here](../../../developers/design-docs/cuda-backend.md/#example). ## Performance diff --git a/docs/website/docs/developers/debugging/compile-time-regressions.md b/docs/website/docs/developers/debugging/compile-time-regressions.md index 53e961a535f2..a6ea76c211bf 100644 --- a/docs/website/docs/developers/debugging/compile-time-regressions.md +++ b/docs/website/docs/developers/debugging/compile-time-regressions.md @@ -44,7 +44,7 @@ Building the compiler from source and using specific commits in IREE, though it typically won't let you step through changes in submodules (e.g. MLIR updates in `third_party/llvm-project/`). -**Tip**: [Configure ccache](../building/cmake_with_ccache.md) if you'll be +**Tip**: [Configure ccache](../building/cmake-with-ccache.md) if you'll be rebuilding the compiler while bisecting A manual workflow with `git bisect` looks like this: @@ -187,7 +187,7 @@ paint a complete picture and requires waiting for compilation to finish. See our documentation on -[profiling with Tracy](../performance/profiling_with_tracy.md). For compile +[profiling with Tracy](../performance/profiling-with-tracy.md). For compile time regressions, pay particular attention to the different compilation phases (Flow/Stream/HAL), how many times `TranslateExecutablesPass` runs, and if there are outlier passes that take significantly longer to run than others. diff --git a/docs/website/docs/developers/debugging/integration-tests.md b/docs/website/docs/developers/debugging/integration-tests.md index 26e719a998a7..5761c48f9a10 100644 --- a/docs/website/docs/developers/debugging/integration-tests.md +++ b/docs/website/docs/developers/debugging/integration-tests.md @@ -10,7 +10,7 @@ Feel free to reach out to @hanhanW or ask questions on Discord for more help. * Models themselves can be large, and IREE breaks models into dispatches/kernels and then launches those individually. Program outputs could diverge starting from any individual launch. To get a smaller reproducer, you can use -[--iree-flow-trace-dispatch-tensors](../general/developer_overview.md#-iree-flow-trace-dispatch-tensors). +[--iree-flow-trace-dispatch-tensors](../general/developer-overview.md#-iree-flow-trace-dispatch-tensors). * You can compare the logs between builds/backends to get an idea about which dispatch results in wrong outputs. The dumped inputs can be reused in a flagfile. diff --git a/docs/website/docs/developers/general/contributing.md b/docs/website/docs/developers/general/contributing.md index 807a734ec0c9..50416303a12f 100644 --- a/docs/website/docs/developers/general/contributing.md +++ b/docs/website/docs/developers/general/contributing.md @@ -90,7 +90,7 @@ this behavior can be disabled with `skip-llvm-integrate-benchmark`. The as part of benchmarking. It accepts a comma-separated list of benchmark presets. This combines with labels added to the PR (which are a more limited set of options). See the -[benchmark suites documentation](../performance/benchmark_suites.md). +[benchmark suites documentation](../performance/benchmark-suites.md). The `runner-env` option controls which runner environment to target for our self-hosted runners. We maintain a test environment to allow testing out new @@ -128,18 +128,18 @@ runs. For example, this PR opted in to running the `build_test_all_windows` job: -![ci-extra](./contributing_ci-extra.png) +![ci-extra](./contributing-ci-extra.png) The enabled jobs can be viewed from the Summary page of an action run: -![ci_enabled_jobs](./contributing_ci_enabled_jobs.png) +![ci_enabled_jobs](./contributing-ci-enabled-jobs.png) ## Contributor tips These are opinionated tips documenting workflows that some members of the team have found useful. They are focused on meta-tooling, not on IREE code specifically (you will find the latter in the -[Developer Overview](./developer_overview.md)). +[Developer Overview](./developer-overview.md)). !!! note @@ -160,7 +160,7 @@ of `git` and GitHub and suggests some specific ways of using it. * VSCode: . The most commonly used IDE amongst IREE developers. * [Ccache](https://ccache.dev/), a fast C/C++ compiler cache. See our - [CMake with `ccache`](../building/cmake_with_ccache.md) page + [CMake with `ccache`](../building/cmake-with-ccache.md) page ### Git structure diff --git a/docs/website/docs/developers/general/developer-overview.md b/docs/website/docs/developers/general/developer-overview.md index 1dd173c68283..94b51f00d08f 100644 --- a/docs/website/docs/developers/general/developer-overview.md +++ b/docs/website/docs/developers/general/developer-overview.md @@ -137,7 +137,7 @@ $ ../iree-build/tools/iree-run-module \ The `iree-check-module` program takes an already translated IREE module as input and executes it as a series of [googletest](https://github.com/google/googletest) tests. This is the test -runner for the IREE [check framework](./testing_guide.md#end-to-end-tests). +runner for the IREE [check framework](./testing-guide.md#end-to-end-tests). ```shell $ ../iree-build/tools/iree-compile \ diff --git a/docs/website/docs/developers/general/release-management.md b/docs/website/docs/developers/general/release-management.md index 7cda3fb07398..33a3a575ec46 100644 --- a/docs/website/docs/developers/general/release-management.md +++ b/docs/website/docs/developers/general/release-management.md @@ -33,6 +33,6 @@ request that some feature make the cut. 2. Open the release on GitHub. Rename the release from "candidate" to "stable", uncheck the option for "pre-release", and check the option for "latest". - ![rename_release](./release_renaming.png) + ![rename_release](./release-renaming.png) - ![promote_release](./release_promotion.png) + ![promote_release](./release-promotion.png) diff --git a/docs/website/docs/developers/general/testing-guide.md b/docs/website/docs/developers/general/testing-guide.md index df26f7363026..59b08e1267d9 100644 --- a/docs/website/docs/developers/general/testing-guide.md +++ b/docs/website/docs/developers/general/testing-guide.md @@ -127,7 +127,7 @@ export CTEST_PARALLEL_LEVEL=$(nproc) To use the Vulkan backend as test driver, you may need to select between a Vulkan implementation from SwiftShader and multiple Vulkan-capable hardware devices. This can be done via environment variables. See the -[generic Vulkan setup](../vulkan_environment_setup.md#useful-environment-variables) +[generic Vulkan setup](../vulkan-environment-setup.md#useful-environment-variables) page for details regarding these variables. For Bazel, you can persist the configuration in `user.bazelrc` to save typing. @@ -226,7 +226,7 @@ To run e2e model tests in [generated_e2e_model_tests.cmake](https://github.com/openxla/iree/tree/main/tests/e2e/stablehlo_models/generated_e2e_model_tests.cmake), because of their dependencies, `-DIREE_BUILD_E2E_TEST_ARTIFACTS=ON` needs to be set when configuring CMake. Also see -[IREE Benchmark Suite Prerequisites](../performance/benchmark_suites.md#prerequisites) +[IREE Benchmark Suite Prerequisites](../performance/benchmark-suites.md#prerequisites) for required packages. ### Running a Test diff --git a/docs/website/docs/developers/performance/benchmark-suites.md b/docs/website/docs/developers/performance/benchmark-suites.md index 0255a3e743da..9c8f56520cb9 100644 --- a/docs/website/docs/developers/performance/benchmark-suites.md +++ b/docs/website/docs/developers/performance/benchmark-suites.md @@ -130,7 +130,7 @@ Note that: - To run x86_64 benchmarks, right now `--cpu_uarch` needs to be provided and only `CascadeLake` is available currently. - To build traced benchmark tools, see - [Profiling with Tracy](profiling_with_tracy.md). + [Profiling with Tracy](profiling-with-tracy.md). Filters can be used to select the benchmarks: diff --git a/docs/website/docs/developers/performance/profiling-cpu-events.md b/docs/website/docs/developers/performance/profiling-cpu-events.md index f3b7cc21f78d..7477490c92c7 100644 --- a/docs/website/docs/developers/performance/profiling-cpu-events.md +++ b/docs/website/docs/developers/performance/profiling-cpu-events.md @@ -87,7 +87,7 @@ report` and `perf annotate`. `perf report` breaks down the event counts by symbol. In the default case where what was sampled was time, this is just an ordinary profile by symbol name, no different than what could be viewed in other profilers such as -[Tracy](profiling_with_tracy.md). Where it gets really interesting is when the +[Tracy](profiling-with-tracy.md). Where it gets really interesting is when the profile was recording a specific event type, as in the above `-e L1-dcache-load-misses` example: @@ -183,7 +183,7 @@ of achieving the same thing. However: * The common case of annotating by time, as opposed to annotating by CPU event, - is supported by [Tracy](profiling_with_tracy.md). + is supported by [Tracy](profiling-with-tracy.md). * Annotating by CPU event is inherently not working due to hardware limitations of the ARM CPUs found in Android devices. That is, the hardware is too imprecise at pinning an event to a particular instruction. diff --git a/docs/website/docs/developers/performance/profiling-gpu-vulkan.md b/docs/website/docs/developers/performance/profiling-gpu-vulkan.md index f5b2861193b9..a92d6d9a11a6 100644 --- a/docs/website/docs/developers/performance/profiling-gpu-vulkan.md +++ b/docs/website/docs/developers/performance/profiling-gpu-vulkan.md @@ -1,6 +1,6 @@ # Profiling GPUs using Vulkan -[Tracy](./profiling_with_tracy.md) offers great insights into CPU/GPU +[Tracy](./profiling-with-tracy.md) offers great insights into CPU/GPU interactions and Vulkan API usage details. However, information at a finer granularity, especially inside a particular shader dispatch, is missing. To supplement general purpose tools diff --git a/docs/website/docs/developers/performance/profiling.md b/docs/website/docs/developers/performance/profiling.md index acc2f2212c7a..a5a42746be63 100644 --- a/docs/website/docs/developers/performance/profiling.md +++ b/docs/website/docs/developers/performance/profiling.md @@ -8,18 +8,18 @@ behavior in more depth, there are various ways to ## Tracy Tracy is a profiler that's been used for a wide range of profiling tasks on -IREE. Refer to [Profiling with Tracy](./profiling_with_tracy.md). +IREE. Refer to [Profiling with Tracy](./profiling-with-tracy.md). ## Vulkan GPU Profiling -[Tracy](./profiling_with_tracy.md) offers great insights into CPU/GPU +[Tracy](./profiling-with-tracy.md) offers great insights into CPU/GPU interactions and Vulkan API usage details. However, information at a finer granularity, especially inside a particular shader dispatch, is missing. To supplement general purpose tools like Tracy, vendor-specific tools can be used. -Refer to [Profiling GPUs using Vulkan](./profiling_gpu_vulkan.md). +Refer to [Profiling GPUs using Vulkan](./profiling-gpu-vulkan.md). ## CPU cache and other CPU event profiling For some advanced CPU profiling needs such as querying CPU cache and other events, one may need to use some OS-specific profilers. See -[profiling_cpu_events.md](./profiling_cpu_events.md). +[Profiling CPUs](./profiling-cpu-events.md). diff --git a/docs/website/overrides/404.html b/docs/website/overrides/404.html index 64dbee56e558..fd3a1299643d 100644 --- a/docs/website/overrides/404.html +++ b/docs/website/overrides/404.html @@ -6,8 +6,6 @@

Sorry, we couldn't find that page.

-

The docs/developers/ directory on GitHub might be helpful. -

Click here to go back to the home page.

{% endblock %} diff --git a/experimental/web/generate_web_metrics.sh b/experimental/web/generate_web_metrics.sh index 1bd8ff7a2504..0bbe785ff679 100644 --- a/experimental/web/generate_web_metrics.sh +++ b/experimental/web/generate_web_metrics.sh @@ -85,7 +85,7 @@ wget -nc https://storage.googleapis.com/iree-model-artifacts/MobileNetV3SmallSta ############################################################################### # Note: you can also download imported programs from CI runs: -# https://github.com/openxla/iree/blob/main/docs/developers/developing_iree/benchmark_suites.md#fetching-benchmark-artifacts-from-ci +# https://iree.dev/developers/performance/benchmark-suites/#fetching-benchmark-artifacts-from-ci IREE_IMPORT_TFLITE_PATH=iree-import-tflite diff --git a/experimental/web/sample_dynamic/index.html b/experimental/web/sample_dynamic/index.html index f95c622b6cc9..cb62dac1d149 100644 --- a/experimental/web/sample_dynamic/index.html +++ b/experimental/web/sample_dynamic/index.html @@ -53,7 +53,7 @@

IREE Dynamic Web Sample

This tool works similarly to iree-run-module - (docs). + (docs).
It loads a compiled IREE program then lets you call exported functions.
Note: Some outputs are logged to the console.

diff --git a/experimental/web/sample_webgpu/index.html b/experimental/web/sample_webgpu/index.html index 81c3709de210..2461e0949d50 100644 --- a/experimental/web/sample_webgpu/index.html +++ b/experimental/web/sample_webgpu/index.html @@ -53,7 +53,7 @@

IREE WebGPU Sample

This tool works similarly to iree-run-module - (docs). + (docs).
It loads a compiled IREE program then lets you call exported functions.
Note: Some outputs are logged to the console.

diff --git a/runtime/src/iree/schemas/bytecode_module_def.fbs b/runtime/src/iree/schemas/bytecode_module_def.fbs index a281e97db286..3ff1d502e51a 100644 --- a/runtime/src/iree/schemas/bytecode_module_def.fbs +++ b/runtime/src/iree/schemas/bytecode_module_def.fbs @@ -42,7 +42,7 @@ table FunctionSignatureDef { // Function-level attributes, if any. // These are typically used to communicate additional ABI metadata needed // for dynamic invocation and host language mapping. - // See: docs/developers/design_docs/function_abi.md + // See: https://iree.dev/developers/design-docs/function-abi/ attrs:[AttrDef]; } diff --git a/runtime/src/iree/vm/module.h b/runtime/src/iree/vm/module.h index 2e393d14eac6..3b7bacf059d3 100644 --- a/runtime/src/iree/vm/module.h +++ b/runtime/src/iree/vm/module.h @@ -539,7 +539,7 @@ IREE_API_EXPORT iree_string_view_t iree_vm_function_lookup_attr_by_name( // of the module. Note that not all functions have reflection attributes. // // For more information on the function ABI and its reflection metadata see: -// docs/developers/design_docs/function_abi.md +// https://iree.dev/developers/design-docs/function-abi/. // // Returns IREE_STATUS_OUT_OF_RANGE if index >= the number of attributes for // the function. diff --git a/samples/custom_module/basic/README.md b/samples/custom_module/basic/README.md index 60522d5a4364..fc414318094f 100644 --- a/samples/custom_module/basic/README.md +++ b/samples/custom_module/basic/README.md @@ -13,7 +13,7 @@ C++ module wrapper layer in [`module.cc`](./module.cc), and called by example in [`main.c`](./main.c). This document uses terminology that can be found in the documentation of -[IREE's execution model](https://github.com/openxla/iree/blob/main/docs/developers/design_docs/execution_model.md). +[IREE's execution model](https://iree.dev/developers/design-docs/invocation-execution-model/). See [IREE's extensibility mechanisms](https://iree.dev/reference/extensions/) documentation for more information specific to extenting IREE and alternative approaches to doing so. diff --git a/tests/e2e/linalg/BUILD.bazel b/tests/e2e/linalg/BUILD.bazel index 60320993b099..6e332f3f5d48 100644 --- a/tests/e2e/linalg/BUILD.bazel +++ b/tests/e2e/linalg/BUILD.bazel @@ -8,7 +8,6 @@ # Each test file should have a name matching the corresponding TOSA op and test only the # functionality of that op (though may make use of other ops where necessary). Tests should be # written using the IREE Check framework. -# See https://github.com/openxla/iree/blob/main/docs/developers/developing_iree/testing_guide.md#iree-core-end-to-end-tests. load("//build_tools/bazel:enforce_glob.bzl", "enforce_glob") load("//build_tools/bazel:iree_check_test.bzl", "iree_check_single_backend_test_suite") diff --git a/tests/e2e/stablehlo_ops/BUILD.bazel b/tests/e2e/stablehlo_ops/BUILD.bazel index 0c29883118e8..adeb3ea5ae21 100644 --- a/tests/e2e/stablehlo_ops/BUILD.bazel +++ b/tests/e2e/stablehlo_ops/BUILD.bazel @@ -8,7 +8,6 @@ # Each test file should have a name matching the corresponding StableHLO op and test only the # functionality of that op (though may make use of other ops where necessary). Tests should be # written using the IREE Check framework and should always pass on the reference VMVX backend. -# See https://github.com/openxla/iree/blob/main/docs/developers/developing_iree/testing_guide.md#iree-core-end-to-end-tests. load("//build_tools/bazel:enforce_glob.bzl", "enforce_glob") load("//build_tools/bazel:iree_check_test.bzl", "iree_check_single_backend_test_suite") diff --git a/tests/e2e/tosa_ops/BUILD.bazel b/tests/e2e/tosa_ops/BUILD.bazel index 37ac8a30667e..4e59a217d191 100644 --- a/tests/e2e/tosa_ops/BUILD.bazel +++ b/tests/e2e/tosa_ops/BUILD.bazel @@ -8,7 +8,6 @@ # Each test file should have a name matching the corresponding TOSA op and test only the # functionality of that op (though may make use of other ops where necessary). Tests should be # written using the IREE Check framework. -# See https://github.com/openxla/iree/blob/main/docs/developers/developing_iree/testing_guide.md#iree-core-end-to-end-tests. load("//build_tools/bazel:enforce_glob.bzl", "enforce_glob") load("//build_tools/bazel:iree_check_test.bzl", "iree_check_single_backend_test_suite") From 09bb1ccfed4b1c02dde93da710e430b9a92e6d39 Mon Sep 17 00:00:00 2001 From: Scott Todd Date: Thu, 2 Nov 2023 13:45:19 -0700 Subject: [PATCH 14/16] Fill in index page. --- docs/website/docs/developers/index.md | 7 +++---- 1 file changed, 3 insertions(+), 4 deletions(-) diff --git a/docs/website/docs/developers/index.md b/docs/website/docs/developers/index.md index eeff4100c7e9..226be22ab4be 100644 --- a/docs/website/docs/developers/index.md +++ b/docs/website/docs/developers/index.md @@ -1,7 +1,6 @@ # Developers -TODO +These pages cover topics useful for project maintainers and contributors. -* Disclaimer about freshness -* Overview for each section -* Difference between this category and the other website categories +!!! caution + Some of these pages may be stale. Contributions are always welcome! From e6735cd4ff1ca59524bc6bc3c91dd7d87ba1a1c6 Mon Sep 17 00:00:00 2001 From: Scott Todd Date: Thu, 2 Nov 2023 13:52:02 -0700 Subject: [PATCH 15/16] Add a few links back with updated URLs. --- tests/e2e/linalg/BUILD.bazel | 1 + tests/e2e/stablehlo_ops/BUILD.bazel | 1 + tests/e2e/tosa_ops/BUILD.bazel | 1 + 3 files changed, 3 insertions(+) diff --git a/tests/e2e/linalg/BUILD.bazel b/tests/e2e/linalg/BUILD.bazel index 6e332f3f5d48..a4a7e5527a1b 100644 --- a/tests/e2e/linalg/BUILD.bazel +++ b/tests/e2e/linalg/BUILD.bazel @@ -8,6 +8,7 @@ # Each test file should have a name matching the corresponding TOSA op and test only the # functionality of that op (though may make use of other ops where necessary). Tests should be # written using the IREE Check framework. +# See https://iree.dev/developers/general/testing-guide/#iree-core-end-to-end-e2e-tests. load("//build_tools/bazel:enforce_glob.bzl", "enforce_glob") load("//build_tools/bazel:iree_check_test.bzl", "iree_check_single_backend_test_suite") diff --git a/tests/e2e/stablehlo_ops/BUILD.bazel b/tests/e2e/stablehlo_ops/BUILD.bazel index adeb3ea5ae21..14852a7a56c3 100644 --- a/tests/e2e/stablehlo_ops/BUILD.bazel +++ b/tests/e2e/stablehlo_ops/BUILD.bazel @@ -8,6 +8,7 @@ # Each test file should have a name matching the corresponding StableHLO op and test only the # functionality of that op (though may make use of other ops where necessary). Tests should be # written using the IREE Check framework and should always pass on the reference VMVX backend. +# See https://iree.dev/developers/general/testing-guide/#iree-core-end-to-end-e2e-tests. load("//build_tools/bazel:enforce_glob.bzl", "enforce_glob") load("//build_tools/bazel:iree_check_test.bzl", "iree_check_single_backend_test_suite") diff --git a/tests/e2e/tosa_ops/BUILD.bazel b/tests/e2e/tosa_ops/BUILD.bazel index 4e59a217d191..a554c080acb8 100644 --- a/tests/e2e/tosa_ops/BUILD.bazel +++ b/tests/e2e/tosa_ops/BUILD.bazel @@ -8,6 +8,7 @@ # Each test file should have a name matching the corresponding TOSA op and test only the # functionality of that op (though may make use of other ops where necessary). Tests should be # written using the IREE Check framework. +# See https://iree.dev/developers/general/testing-guide/#iree-core-end-to-end-e2e-tests. load("//build_tools/bazel:enforce_glob.bzl", "enforce_glob") load("//build_tools/bazel:iree_check_test.bzl", "iree_check_single_backend_test_suite") From 4327c90685556aed4c0222c4ac103646adbcc137 Mon Sep 17 00:00:00 2001 From: Scott Todd Date: Thu, 2 Nov 2023 14:48:31 -0700 Subject: [PATCH 16/16] Add more prominent link to website in website source README. --- docs/website/README.md | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/docs/website/README.md b/docs/website/README.md index bf13a8890526..84cd5e8b1d58 100644 --- a/docs/website/README.md +++ b/docs/website/README.md @@ -1,10 +1,10 @@ # IREE User-Facing Documentation Website -This directory contains the source and assets for IREE's website, hosted on -[GitHub Pages](https://pages.github.com/). +This directory contains the source and assets for . The website is generated using [MkDocs](https://www.mkdocs.org/), with the -[Material for MkDocs](https://squidfunk.github.io/mkdocs-material/) theme. +[Material for MkDocs](https://squidfunk.github.io/mkdocs-material/) theme and +is served using [GitHub Pages](https://pages.github.com/). ## How to edit this documentation