Merge remote-tracking branch 'origin/main' into bazel-gpu
diff --git a/.github/workflows/publish_docs.yml b/.github/workflows/publish_docs.yml index 123f5a4..457da97 100644 --- a/.github/workflows/publish_docs.yml +++ b/.github/workflows/publish_docs.yml
@@ -33,9 +33,6 @@ uses: actions/checkout@v2 with: token: ${{ secrets.GITHUB_WRITE_ACCESS_TOKEN }} - - name: Fetching gh-pages branch - run: | - git fetch origin gh-pages - name: Initializing submodules run: ./scripts/git/submodule_versions.py init - name: Installing Ninja build @@ -45,6 +42,9 @@ ./build_tools/cmake/build_docs.sh # Patch the MarkDown files with front matter for rendering ./scripts/prepare_doc_publication.py ${IREE_DOC_BUILD_DIR}/doc + - name: Fetching gh-pages branch + run: | + git fetch origin gh-pages - name: Updating gh-pages branch run: | git checkout -f gh-pages
diff --git a/.github/workflows/update_tf.yml b/.github/workflows/update_tf.yml index 86c55d9..b116f99 100644 --- a/.github/workflows/update_tf.yml +++ b/.github/workflows/update_tf.yml
@@ -54,6 +54,4 @@ Automated submodule bump from .github/workflows/update_tf.yml committer: "Submodule Update Action <iree-github-actions-bot@google.com>" - # TODO(gcmn): Figure out a way to assign this to someone dynamically. - reviewers: gmngeoffrey branch: "auto_submodule_update"
diff --git a/.gitmodules b/.gitmodules index 13d12e9..3bde2de 100644 --- a/.gitmodules +++ b/.gitmodules
@@ -50,3 +50,6 @@ [submodule "third_party/marl"] path = third_party/marl url = https://github.com/google/marl.git +[submodule "third_party/flatcc"] + path = third_party/flatcc + url = https://github.com/dvidelabs/flatcc.git
diff --git a/.style.yapf b/.style.yapf new file mode 100644 index 0000000..9ef1dc1 --- /dev/null +++ b/.style.yapf
@@ -0,0 +1,4 @@ +[style] + based_on_style = google + column_limit = 80 + indent_width = 2
diff --git a/CMakeLists.txt b/CMakeLists.txt index 22bb32f..155f08d 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt
@@ -93,6 +93,10 @@ if( IREE_HAL_DRIVERS_TO_BUILD STREQUAL "all" ) set( IREE_HAL_DRIVERS_TO_BUILD ${IREE_ALL_HAL_DRIVERS} ) + # For cross compilation towords Android, we don't want LLVM JIT HAL driver. + if(ANDROID) + list(REMOVE_ITEM IREE_HAL_DRIVERS_TO_BUILD LLVM) + endif() endif() message(STATUS "Building HAL drivers ${IREE_HAL_DRIVERS_TO_BUILD}") @@ -112,8 +116,8 @@ # List of all target backends to be built by default: set(IREE_ALL_TARGET_BACKENDS # TODO(scotttodd): LLVMAOT - LLVMIR - Vulkan_SPIRV + LLVM-IR + Vulkan-SPIRV VMLA ) @@ -201,10 +205,7 @@ include(iree_tablegen_doc) include(iree_cc_embed_data) include(iree_bytecode_module) -include(iree_pybind_cc_library) -include(iree_py_extension) -include(iree_py_library) -include(iree_py_test) +include(iree_multipy) include(iree_lit_test) include(iree_add_all_subdirs) include(iree_check_test) @@ -251,7 +252,6 @@ # dependency is added prior to including this configuration. #------------------------------------------------------------------------------- - # Adds bundled projects that must be included after the LLVM directory has # been added and within the scope of its settings (i.e. build type override, # etc). @@ -314,7 +314,6 @@ add_bundled_mlir_dependent_projects() endif() - #------------------------------------------------------------------------------- # Non-LLVM Dependencies #------------------------------------------------------------------------------- @@ -329,7 +328,8 @@ find_package(PythonInterp 3 REQUIRED) endif() if(${IREE_BUILD_PYTHON_BINDINGS}) - find_package(PythonLibs 3 REQUIRED) + # Note: Optional because python libs can be manually specified. + find_package(PythonLibs 3) endif() list(APPEND CMAKE_MODULE_PATH @@ -337,20 +337,24 @@ ) include(external_cc_library) +include(flatbuffer_c_library) include(flatbuffer_cc_library) +add_subdirectory(build_tools/third_party/flatcc EXCLUDE_FROM_ALL) +add_subdirectory(build_tools/third_party/renderdoc_api EXCLUDE_FROM_ALL) add_subdirectory(build_tools/third_party/ruy EXCLUDE_FROM_ALL) add_subdirectory(third_party/googletest EXCLUDE_FROM_ALL) add_subdirectory(third_party/abseil-cpp EXCLUDE_FROM_ALL) add_subdirectory(third_party/flatbuffers EXCLUDE_FROM_ALL) +add_subdirectory(third_party/flatcc EXCLUDE_FROM_ALL) add_subdirectory(third_party/vulkan_headers EXCLUDE_FROM_ALL) -add_subdirectory(build_tools/third_party/renderdoc_api EXCLUDE_FROM_ALL) if(CMAKE_CROSSCOMPILING) # We need flatc to generate some source code. When cross-compiling, we need # to make sure the flatc binary is configured under host environment. iree_declare_host_excutable(flatc BUILDONLY) + iree_declare_host_excutable(flatcc BUILDONLY) # Set the FLATBUFFERS_FLATC_EXECUTABLE. It controls where to find the flatc # binary in BuildFlatBuffers(). @@ -358,12 +362,21 @@ # Add a custom target to copy the flatc to the binary directory. add_custom_target(iree_host_flatc - COMMAND "${CMAKE_COMMAND}" -E copy_if_different - "${IREE_HOST_BINARY_ROOT}/third_party/flatbuffers/flatc${IREE_HOST_EXECUTABLE_SUFFIX}" - "${IREE_HOST_BINARY_ROOT}/bin" + COMMAND + "${CMAKE_COMMAND}" -E copy_if_different + "${IREE_HOST_BINARY_ROOT}/third_party/flatbuffers/flatc${IREE_HOST_EXECUTABLE_SUFFIX}" + "${IREE_HOST_BINARY_ROOT}/bin" DEPENDS iree_host_build_flatc COMMENT "Installing host flatc..." ) + add_custom_target(iree_host_flatcc + COMMAND + "${CMAKE_COMMAND}" -E copy_if_different + "${IREE_HOST_BINARY_ROOT}/third_party/flatcc/flatcc${IREE_HOST_EXECUTABLE_SUFFIX}" + "${IREE_HOST_BINARY_ROOT}/bin" + DEPENDS iree_host_build_flatcc + COMMENT "Installing host flatcc..." + ) endif() if(${IREE_BUILD_COMPILER}) @@ -387,7 +400,10 @@ endif() if(${IREE_BUILD_PYTHON_BINDINGS}) + # NOTE: The multipy defaults come from pybind's configuration and must come + # after. This should be pulled in locally at some point. add_subdirectory(third_party/pybind11 EXCLUDE_FROM_ALL) + iree_multipy_configure() endif() #-------------------------------------------------------------------------------
diff --git a/CONTRIBUTING.md b/CONTRIBUTING.md index a906067..0f7897d 100644 --- a/CONTRIBUTING.md +++ b/CONTRIBUTING.md
@@ -58,7 +58,7 @@ ## Peculiarities Our documentation on -[repository management](https://github.com/google/iree/blob/main/docs/repository_management.md) +[repository management](https://github.com/google/iree/blob/main/docs/developing_iree/repository_management.md) has more information on some of the oddities in our repository setup and workflows. For the most part, these should be transparent to normal developer workflows.
diff --git a/README.md b/README.md index e6de0e2..b1fd668 100644 --- a/README.md +++ b/README.md
@@ -37,8 +37,8 @@ working on enabling macOS support. For deployment, IREE aims to additionally cover Android and iOS. -Please see the [Getting Started](https://google.github.io/iree/GetStarted) pages -on IREE's [documentation hub](https://google.github.io/iree) to configure, +Please see the [Getting Started](https://google.github.io/iree/get_started) +pages on IREE's [documentation hub](https://google.github.io/iree) to configure, compile, and run IREE in your favorite development environment! ## Documentation and Talks @@ -68,7 +68,7 @@ The architecture of IREE is best illustrated by the following picture: - + Being compilation-based means IREE does not have a traditional runtime that dispatches "ops" to their fat kernel implementations. What IREE provides is a @@ -100,8 +100,8 @@ ## Roadmap and Milestones IREE is still at its early stage; we have lots of exciting future plans. Please -check out the [long-term design roadmap](./docs/roadmap_design.md) and -[short-term focus areas](./docs/roadmap.md). +check out the [long-term design roadmap](./docs/design_roadmap.md) and +[short-term focus areas](./docs/milestones.md). We use [GitHub Projects](https://github.com/google/iree/projects) to track various IREE components and
diff --git a/SUBMODULE_VERSIONS b/SUBMODULE_VERSIONS index 1e315f6..b802f7d 100644 --- a/SUBMODULE_VERSIONS +++ b/SUBMODULE_VERSIONS
@@ -2,8 +2,9 @@ daff5fead3fbe22c6fc58310ca3f49caf117f185 third_party/benchmark 4c13807b7d43ff0946b7ffea0ae3aee9e611d778 third_party/dear_imgui a5d9d0f7d368054fd1691aedf1db4116efcc233e third_party/flatbuffers +4fb0ff7069bd88ee85902f4d0bb62794e5f6d021 third_party/flatcc f2fb48c3b3d79a75a88a99fba6576b25d42ec528 third_party/googletest -de0c6bd56b41081f1b89a1c7a0bf2597fd6d0104 third_party/llvm-project +99ad956fdaee5398fdcf46fa49cb433cf52dc461 third_party/llvm-project 17b12a4481daa150e2d1ea3ada086b551b856707 third_party/marl 67f3ccebee84f3488b46a8d3ac005178c52ff264 third_party/mlir-emitc 80d452484c5409444b0ec19383faa84bb7a4d351 third_party/pybind11 @@ -11,7 +12,7 @@ b73f111094da3e380a1774b56b15f16c90ae8e23 third_party/sdl2 f8bf11a0253a32375c32cad92c841237b96696c0 third_party/spirv_headers 57eb48aed36160c4876bc8310d9ca84d42ee9e2a third_party/swiftshader -e36aca0132fbcde0bc820d56185e3078f97a879d third_party/tensorflow +e4a48da690fac3443825b535ec02cb31c5625337 third_party/tensorflow 864d86e8b6d21449474db5e9313dbff90aa9c24f third_party/tracy 9bd3f561bcee3f01d22912de10bb07ce4e23d378 third_party/vulkan_headers 909f36b714c9239ee0b112a321220213a474ba53 third_party/vulkan_memory_allocator
diff --git a/WORKSPACE b/WORKSPACE index 507a82a..cf098ce 100644 --- a/WORKSPACE +++ b/WORKSPACE
@@ -63,7 +63,7 @@ rbe_autoconfig( name = "rbe_default", base_container_digest = "sha256:1a8ed713f40267bb51fe17de012fa631a20c52df818ccb317aaed2ee068dfc61", - digest = "sha256:b59d8cc422b03524394d4d05e443bf38d4fe96fab06197b34174de01572e8161", + digest = "sha256:bc2d61ad05453928e67b434ae019e7d050dda46c091270f2b81b2f09da2276ce", registry = "gcr.io", repository = "iree-oss/rbe-toolchain", use_checked_in_confs = "Force", @@ -183,6 +183,13 @@ path = "third_party/flatbuffers", ) +maybe( + new_local_repository, + name = "com_github_dvidelabs_flatcc", + build_file = "build_tools/third_party/flatcc/BUILD.overlay", + path = "third_party/flatcc", +) + # TODO(scotttodd): TensorFlow is squatting on the vulkan_headers repo name, so # we use a temporary one until resolved. Theirs is set to an outdated version. maybe(
diff --git a/bindings/java/com/google/iree/Context.java b/bindings/java/com/google/iree/Context.java index b2a4399..2b42c79 100644 --- a/bindings/java/com/google/iree/Context.java +++ b/bindings/java/com/google/iree/Context.java
@@ -54,6 +54,15 @@ } } + public Function resolveFunction(String name) throws Exception { + Function function = new Function(); + Status status = Status.fromCode(nativeResolveFunction(function.getNativeAddress(), name)); + if (!status.isOk()) { + throw status.toException("Could not resolve function"); + } + return function; + } + public int getId() { return nativeGetId(); } @@ -82,6 +91,8 @@ private native int nativeRegisterModules(long[] moduleAddresses); + private native int nativeResolveFunction(long functionAddress, String name); + private native void nativeFree(); private native int nativeGetId();
diff --git a/bindings/java/com/google/iree/Function.java b/bindings/java/com/google/iree/Function.java new file mode 100644 index 0000000..bf9c02a --- /dev/null +++ b/bindings/java/com/google/iree/Function.java
@@ -0,0 +1,38 @@ +/* + * Copyright 2020 Google LLC + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +package com.google.iree; + +/** A function reference. */ +final class Function { + public Function() { + nativeAddress = nativeNew(); + } + + public long getNativeAddress() { + return nativeAddress; + } + + public void free() { + nativeFree(); + } + + private final long nativeAddress; + + private native long nativeNew(); + + private native void nativeFree(); +}
diff --git a/bindings/java/com/google/iree/native/context_jni.cc b/bindings/java/com/google/iree/native/context_jni.cc index 341dad2..de11e74 100644 --- a/bindings/java/com/google/iree/native/context_jni.cc +++ b/bindings/java/com/google/iree/native/context_jni.cc
@@ -15,6 +15,7 @@ #include <jni.h> #include "bindings/java/com/google/iree/native/context_wrapper.h" +#include "bindings/java/com/google/iree/native/function_wrapper.h" #include "bindings/java/com/google/iree/native/instance_wrapper.h" #include "bindings/java/com/google/iree/native/module_wrapper.h" #include "iree/base/logging.h" @@ -23,6 +24,7 @@ #define JNI_PREFIX(METHOD) Java_com_google_iree_Context_##METHOD using iree::java::ContextWrapper; +using iree::java::FunctionWrapper; using iree::java::InstanceWrapper; using iree::java::ModuleWrapper; @@ -101,6 +103,21 @@ return (jint)status.code(); } +JNI_FUNC jint JNI_PREFIX(nativeResolveFunction)(JNIEnv* env, jobject thiz, + jlong functionAddress, + jstring name) { + ContextWrapper* context = GetContextWrapper(env, thiz); + CHECK_NE(context, nullptr); + + auto function = (FunctionWrapper*)functionAddress; + const char* native_name = env->GetStringUTFChars(name, /*isCopy=*/nullptr); + + auto status = context->ResolveFunction( + *function, iree_string_view_t{native_name, strlen(native_name)}); + env->ReleaseStringUTFChars(name, native_name); + return (jint)status.code(); +} + JNI_FUNC jint JNI_PREFIX(nativeGetId)(JNIEnv* env, jobject thiz) { ContextWrapper* context = GetContextWrapper(env, thiz); CHECK_NE(context, nullptr);
diff --git a/bindings/java/com/google/iree/native/context_wrapper.cc b/bindings/java/com/google/iree/native/context_wrapper.cc index eb3d354..f329644 100644 --- a/bindings/java/com/google/iree/native/context_wrapper.cc +++ b/bindings/java/com/google/iree/native/context_wrapper.cc
@@ -70,6 +70,13 @@ IREE_LOC); } +Status ContextWrapper::ResolveFunction(const FunctionWrapper& function_wrapper, + iree_string_view_t name) { + return FromApiStatus(iree_vm_context_resolve_function( + context_, name, function_wrapper.function()), + IREE_LOC); +} + int ContextWrapper::id() const { return iree_vm_context_id(context_); } ContextWrapper::~ContextWrapper() {
diff --git a/bindings/java/com/google/iree/native/context_wrapper.h b/bindings/java/com/google/iree/native/context_wrapper.h index 380d587..3924a8c 100644 --- a/bindings/java/com/google/iree/native/context_wrapper.h +++ b/bindings/java/com/google/iree/native/context_wrapper.h
@@ -15,6 +15,7 @@ #ifndef IREE_BINDINGS_JAVA_COM_GOOGLE_IREE_NATIVE_CONTEXT_WRAPPER_H_ #define IREE_BINDINGS_JAVA_COM_GOOGLE_IREE_NATIVE_CONTEXT_WRAPPER_H_ +#include "bindings/java/com/google/iree/native/function_wrapper.h" #include "bindings/java/com/google/iree/native/instance_wrapper.h" #include "bindings/java/com/google/iree/native/module_wrapper.h" #include "iree/base/status.h" @@ -34,6 +35,9 @@ Status RegisterModules(const std::vector<ModuleWrapper*>& module_wrappers); + Status ResolveFunction(const FunctionWrapper& function_wrapper, + iree_string_view_t name); + int id() const; ~ContextWrapper();
diff --git a/bindings/java/com/google/iree/native/function_jni.cc b/bindings/java/com/google/iree/native/function_jni.cc new file mode 100644 index 0000000..5b9542f --- /dev/null +++ b/bindings/java/com/google/iree/native/function_jni.cc
@@ -0,0 +1,49 @@ +// Copyright 2020 Google LLC +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// https://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#include <jni.h> + +#include "bindings/java/com/google/iree/native/function_wrapper.h" +#include "iree/base/logging.h" + +#define JNI_FUNC extern "C" JNIEXPORT +#define JNI_PREFIX(METHOD) Java_com_google_iree_Function_##METHOD + +using iree::java::FunctionWrapper; + +namespace { + +// Returns a pointer to the native IREE function stored by the FunctionWrapper +// object. +static FunctionWrapper* GetFunctionWrapper(JNIEnv* env, jobject obj) { + jclass clazz = env->GetObjectClass(obj); + CHECK(clazz); + + jfieldID field = env->GetFieldID(clazz, "nativeAddress", "J"); + CHECK(field); + + return reinterpret_cast<FunctionWrapper*>(env->GetLongField(obj, field)); +} + +} // namespace + +JNI_FUNC jlong JNI_PREFIX(nativeNew)(JNIEnv* env, jobject thiz) { + return reinterpret_cast<jlong>(new FunctionWrapper()); +} + +JNI_FUNC void JNI_PREFIX(nativeFree)(JNIEnv* env, jobject thiz) { + FunctionWrapper* function = GetFunctionWrapper(env, thiz); + CHECK_NE(function, nullptr); + delete function; +}
diff --git a/bindings/java/com/google/iree/native/function_wrapper.cc b/bindings/java/com/google/iree/native/function_wrapper.cc new file mode 100644 index 0000000..2dfcd62 --- /dev/null +++ b/bindings/java/com/google/iree/native/function_wrapper.cc
@@ -0,0 +1,25 @@ +// Copyright 2020 Google LLC +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// https://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#include "bindings/java/com/google/iree/native/function_wrapper.h" + +namespace iree { +namespace java { + +iree_vm_function_t* FunctionWrapper::function() const { + return function_.get(); +} + +} // namespace java +} // namespace iree
diff --git a/bindings/java/com/google/iree/native/function_wrapper.h b/bindings/java/com/google/iree/native/function_wrapper.h new file mode 100644 index 0000000..8c71dc0 --- /dev/null +++ b/bindings/java/com/google/iree/native/function_wrapper.h
@@ -0,0 +1,37 @@ +// Copyright 2020 Google LLC +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// https://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#ifndef IREE_BINDINGS_JAVA_COM_GOOGLE_IREE_NATIVE_FUNCTION_WRAPPER_H_ +#define IREE_BINDINGS_JAVA_COM_GOOGLE_IREE_NATIVE_FUNCTION_WRAPPER_H_ + +#include <memory> + +#include "iree/vm/module.h" + +namespace iree { +namespace java { + +class FunctionWrapper { + public: + iree_vm_function_t* function() const; + + private: + std::unique_ptr<iree_vm_function_t> function_ = + std::make_unique<iree_vm_function_t>(); +}; + +} // namespace java +} // namespace iree + +#endif // IREE_BINDINGS_JAVA_COM_GOOGLE_IREE_NATIVE_FUNCTION_WRAPPER_H_
diff --git a/bindings/javatests/com/google/iree/IntegrationTest.java b/bindings/javatests/com/google/iree/IntegrationTest.java index 499a9c8..13f3fbb 100644 --- a/bindings/javatests/com/google/iree/IntegrationTest.java +++ b/bindings/javatests/com/google/iree/IntegrationTest.java
@@ -60,6 +60,12 @@ assertNotEquals(ireeContext.getId(), -1); + String functionName = "module.simple_mul"; + Function function = ireeContext.resolveFunction(functionName); + + // TODO(jennik): Invoke the function. + + function.free(); module.free(); ireeContext.free(); instance.free(); @@ -83,6 +89,10 @@ modules.add(module); ireeContext.registerModules(modules); + String functionName = "module.simple_mul"; + Function function = ireeContext.resolveFunction(functionName); + + function.free(); module.free(); ireeContext.free(); instance.free();
diff --git a/bindings/python/build_tools/cmake/iree_py_extension.cmake b/bindings/python/build_tools/cmake/iree_py_extension.cmake deleted file mode 100644 index f3c36b8..0000000 --- a/bindings/python/build_tools/cmake/iree_py_extension.cmake +++ /dev/null
@@ -1,259 +0,0 @@ -# Copyright 2020 Google LLC -# -# Licensed under the Apache License, Version 2.0 (the "License"); -# you may not use this file except in compliance with the License. -# You may obtain a copy of the License at -# -# https://www.apache.org/licenses/LICENSE-2.0 -# -# Unless required by applicable law or agreed to in writing, software -# distributed under the License is distributed on an "AS IS" BASIS, -# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -# See the License for the specific language governing permissions and -# limitations under the License. - -include(CMakeParseArguments) - -if (NOT DEFINED _IREE_PY_EXTENSION_NAMES) - set(_IREE_PY_EXTENSION_NAMES "") -endif() - -# iree_py_extension() -# -# CMake function to imitate Bazel's iree_py_extension rule. -# -# Parameters: -# NAME: name of target -# HDRS: List of public header files for the library -# SRCS: List of source files for the library -# DEPS: List of other libraries to be linked in to the py extension targets -# COPTS: List of private compile options -# DEFINES: List of public defines -# INCLUDES: Include directories to add to dependencies -# LINKOPTS: List of link options -# TYPE: Type of library to be crated: either "MODULE", "SHARED" or "STATIC" (default). -# TESTONLY: When added, this target will only be built if user passes -DIREE_BUILD_TESTS=ON to CMake. - -function(iree_py_extension) - - cmake_parse_arguments( - _RULE - "TESTONLY" - "NAME" - "HDRS;SRCS;COPTS;DEFINES;LINKOPTS;DEPS;INCLUDES;TYPE" - ${ARGN} - ) - - iree_package_ns(_PACKAGE_NS) - # Replace dependencies passed by ::name with ::iree::package::name - list(TRANSFORM _RULE_DEPS REPLACE "^::" "${_PACKAGE_NS}::") - - if(NOT _RULE_TESTONLY OR IREE_BUILD_TESTS) - # Prefix the library with the package name, so we get: iree_package_name. - iree_package_name(_PACKAGE_NAME) - set(_NAME "${_PACKAGE_NAME}_${_RULE_NAME}") - - if(NOT _RULE_TYPE) - set(_RULE_TYPE "STATIC") - endif() - - string(TOUPPER "${_RULE_TYPE}" _uppercase_RULE_TYPE) - if(NOT _uppercase_RULE_TYPE MATCHES "^(STATIC|SHARED|MODULE)") - message(FATAL_ERROR "Unsported library TYPE for iree_pybind_cc_library: ${_RULE_TYPE}") - endif() - - - add_library(${_NAME} ${_uppercase_RULE_TYPE} "") - target_sources(${_NAME} - PRIVATE - ${_RULE_SRCS} - ${_RULE_TEXTUAL_HDRS} - ${_RULE_HDRS} - ) - - target_include_directories(${_NAME} - PUBLIC - "$<BUILD_INTERFACE:${IREE_COMMON_INCLUDE_DIRS}>" - "$<BUILD_INTERFACE:${_RULE_INCLUDES}>" - PRIVATE - ${PYBIND11_INCLUDE_DIR} - ${PYTHON_INCLUDE_DIR} - ) - - target_compile_options(${_NAME} - PRIVATE - ${_RULE_COPTS} - ${IREE_DEFAULT_COPTS} - ) - - target_compile_definitions(${_NAME} - PUBLIC - ${_RULE_DEFINES} - ) - - set_target_properties(${_NAME} PROPERTIES OUTPUT_NAME "${_RULE_NAME}") - - if(NOT _uppercase_RULE_TYPE MATCHES "STATIC") - set_property(TARGET ${_NAME} PROPERTY PREFIX "${PYTHON_MODULE_PREFIX}") - set_property(TARGET ${_NAME} PROPERTY SUFFIX "${PYTHON_MODULE_EXTENSION}") - endif() - - # Alias the iree_package_name library to iree::package::name. - # This lets us more clearly map to Bazel and makes it possible to - # disambiguate the underscores in paths vs. the separators. - add_library(${_PACKAGE_NS}::${_RULE_NAME} ALIAS ${_NAME}) - iree_package_dir(_PACKAGE_DIR) - if(${_RULE_NAME} STREQUAL ${_PACKAGE_DIR}) - # If the library name matches the package then treat it as a default. - # For example, foo/bar/ library 'bar' would end up as 'foo::bar'. - add_library(${_PACKAGE_NS} ALIAS ${_NAME}) - endif() - # Defer computing transitive dependencies and calling target_link_libraries() - # until all libraries have been declared. - # Track target and deps, use in iree_complete_py_extension_link_options() later. - set_property(GLOBAL APPEND PROPERTY _IREE_PY_EXTENSION_NAMES "${_NAME}") - set_property(TARGET ${_NAME} PROPERTY DIRECT_DEPS ${_RULE_DEPS}) - endif() -endfunction() - -# Lists all transitive dependencies of DIRECT_DEPS in TRANSITIVE_DEPS. -function(_iree_transitive_dependencies DIRECT_DEPS TRANSITIVE_DEPS) - set(_TRANSITIVE "") - - foreach(_DEP ${DIRECT_DEPS}) - _iree_transitive_dependencies_helper(${_DEP} _TRANSITIVE) - endforeach(_DEP) - - set(${TRANSITIVE_DEPS} "${_TRANSITIVE}" PARENT_SCOPE) -endfunction() - -# Recursive helper function for _iree_transitive_dependencies. -# Performs a depth-first search through the dependency graph, appending all -# dependencies of TARGET to the TRANSITIVE_DEPS list. -function(_iree_transitive_dependencies_helper TARGET TRANSITIVE_DEPS) - if (NOT TARGET "${TARGET}") - # Excluded from the project, or invalid name? Just ignore. - return() - endif() - - # Resolve aliases, canonicalize name formatting. - get_target_property(_ALIASED_TARGET ${TARGET} ALIASED_TARGET) - if(_ALIASED_TARGET) - set(_TARGET_NAME ${_ALIASED_TARGET}) - else() - string(REPLACE "::" "_" _TARGET_NAME ${TARGET}) - endif() - - set(_RESULT "${${TRANSITIVE_DEPS}}") - if (${_TARGET_NAME} IN_LIST _RESULT) - # Already visited, ignore. - return() - endif() - - # Append this target to the list. Dependencies of this target will be added - # (if valid and not already visited) in recursive function calls. - list(APPEND _RESULT ${_TARGET_NAME}) - - # Check for non-target identifiers again after resolving the alias. - if (NOT TARGET ${_TARGET_NAME}) - return() - endif() - - # Get the list of direct dependencies for this target. - get_target_property(_TARGET_TYPE ${_TARGET_NAME} TYPE) - if(NOT ${_TARGET_TYPE} STREQUAL "INTERFACE_LIBRARY") - get_target_property(_TARGET_DEPS ${_TARGET_NAME} LINK_LIBRARIES) - else() - get_target_property(_TARGET_DEPS ${_TARGET_NAME} INTERFACE_LINK_LIBRARIES) - endif() - - if(_TARGET_DEPS) - # Recurse on each dependency. - foreach(_TARGET_DEP ${_TARGET_DEPS}) - _iree_transitive_dependencies_helper(${_TARGET_DEP} _RESULT) - endforeach(_TARGET_DEP) - endif() - - # Propagate the augmented list up to the parent scope. - set(${TRANSITIVE_DEPS} "${_RESULT}" PARENT_SCOPE) -endfunction() - -# Sets target_link_libraries() on all registered py extensions. -# This must be called after all libraries have been declared. -function(iree_complete_py_extension_link_options) - get_property(_NAMES GLOBAL PROPERTY _IREE_PY_EXTENSION_NAMES) - - foreach(_NAME ${_NAMES}) - get_target_property(_DIRECT_DEPS ${_NAME} DIRECT_DEPS) - - # List all dependencies, including transitive dependencies, then split the - # dependency list into one for whole archive (ALWAYSLINK) and one for - # standard linking (which only links in symbols that are directly used). - _iree_transitive_dependencies("${_DIRECT_DEPS}" _TRANSITIVE_DEPS) - set(_ALWAYS_LINK_DEPS "") - set(_STANDARD_DEPS "") - foreach(_DEP ${_TRANSITIVE_DEPS}) - # Check if _DEP is a library with the ALWAYSLINK property set. - set(_DEP_IS_ALWAYSLINK OFF) - if (TARGET ${_DEP}) - get_target_property(_DEP_TYPE ${_DEP} TYPE) - if(${_DEP_TYPE} STREQUAL "INTERFACE_LIBRARY") - # Can't be ALWAYSLINK since it's an INTERFACE library. - # We also can't even query for the property, since it isn't allowlisted. - else() - get_target_property(_DEP_IS_ALWAYSLINK ${_DEP} ALWAYSLINK) - endif() - endif() - - # Append to the corresponding list of deps. - if(_DEP_IS_ALWAYSLINK) - list(APPEND _ALWAYS_LINK_DEPS ${_DEP}) - - # For MSVC, also add a `-WHOLEARCHIVE:` version of the dep. - # CMake treats -WHOLEARCHIVE[:lib] as a link flag and will not actually - # try to link the library in, so we need the flag *and* the dependency. - # For macOS, also add a `-Wl,-force_load` version of the dep. - if(MSVC) - get_target_property(_ALIASED_TARGET ${_DEP} ALIASED_TARGET) - if (_ALIASED_TARGET) - list(APPEND _ALWAYS_LINK_DEPS "-WHOLEARCHIVE:${_ALIASED_TARGET}") - else() - list(APPEND _ALWAYS_LINK_DEPS "-WHOLEARCHIVE:${_DEP}") - endif() - elseif(APPLE) - get_target_property(_ALIASED_TARGET ${_DEP} ALIASED_TARGET) - if (_ALIASED_TARGET) - list(APPEND _ALWAYS_LINK_DEPS "-Wl,-force_load $<TARGET_FILE:${_ALIASED_TARGET}>") - else() - list(APPEND _ALWAYS_LINK_DEPS "-Wl,-force_load $<TARGET_FILE:${_DEP}>") - endif() - endif() - else() - list(APPEND _STANDARD_DEPS ${_DEP}) - endif() - endforeach(_DEP) - - # Call into target_link_libraries with the lists of deps. - if(MSVC OR APPLE) - target_link_libraries(${_NAME} - PUBLIC - ${_ALWAYS_LINK_DEPS} - ${_STANDARD_DEPS} - PRIVATE - ${_RULE_LINKOPTS} - ${PYTHON_LIBRARY} - ) - else() - target_link_libraries(${_NAME} - PUBLIC - "-Wl,--whole-archive" - ${_ALWAYS_LINK_DEPS} - "-Wl,--no-whole-archive" - ${_STANDARD_DEPS} - PRIVATE - ${_RULE_LINKOPTS} - ${PYTHON_LIBRARY} - ) - endif() - endforeach(_NAME) -endfunction()
diff --git a/bindings/python/build_tools/cmake/iree_py_library.cmake b/bindings/python/build_tools/cmake/iree_py_library.cmake deleted file mode 100644 index 8bc8fdf..0000000 --- a/bindings/python/build_tools/cmake/iree_py_library.cmake +++ /dev/null
@@ -1,51 +0,0 @@ -# Copyright 2020 Google LLC -# -# Licensed under the Apache License, Version 2.0 (the "License"); -# you may not use this file except in compliance with the License. -# You may obtain a copy of the License at -# -# https://www.apache.org/licenses/LICENSE-2.0 -# -# Unless required by applicable law or agreed to in writing, software -# distributed under the License is distributed on an "AS IS" BASIS, -# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -# See the License for the specific language governing permissions and -# limitations under the License. - -include(CMakeParseArguments) - -# iree_py_library() -# -# CMake function to imitate Bazel's iree_py_library rule. -# -# Parameters: -# NAME: name of target -# SRCS: List of source files for the library -# DEPS: List of other targets the test python libraries require - -function(iree_py_library) - - cmake_parse_arguments( - _RULE - "" - "NAME" - "SRCS;DEPS" - ${ARGN} - ) - - iree_package_ns(_PACKAGE_NS) - # Replace dependencies passed by ::name with ::iree::package::name - list(TRANSFORM _RULE_DEPS REPLACE "^::" "${_PACKAGE_NS}::") - - iree_package_name(_PACKAGE_NAME) - set(_NAME "${_PACKAGE_NAME}_${_RULE_NAME}") - - # Add path to each source file - list(TRANSFORM _RULE_SRCS PREPEND "${CMAKE_CURRENT_SOURCE_DIR}/") - - add_custom_target(${_NAME} ALL - COMMAND ${CMAKE_COMMAND} -E copy "${_RULE_SRCS}" "${CMAKE_CURRENT_BINARY_DIR}/" - DEPENDS ${_RULE_DEPS} - ) - -endfunction()
diff --git a/bindings/python/build_tools/cmake/iree_py_test.cmake b/bindings/python/build_tools/cmake/iree_py_test.cmake deleted file mode 100644 index 0fbb6f0..0000000 --- a/bindings/python/build_tools/cmake/iree_py_test.cmake +++ /dev/null
@@ -1,60 +0,0 @@ -# Copyright 2020 Google LLC -# -# Licensed under the Apache License, Version 2.0 (the "License"); -# you may not use this file except in compliance with the License. -# You may obtain a copy of the License at -# -# https://www.apache.org/licenses/LICENSE-2.0 -# -# Unless required by applicable law or agreed to in writing, software -# distributed under the License is distributed on an "AS IS" BASIS, -# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -# See the License for the specific language governing permissions and -# limitations under the License. - -include(CMakeParseArguments) - -# iree_py_test() -# -# CMake function to imitate Bazel's iree_py_test rule. -# -# Parameters: -# NAME: name of test -# SRCS: List of source file -# DEPS: List of deps the test requires -# LABELS: Additional labels to apply to the test. The package path is added -# automatically. - -function(iree_py_test) - if(NOT IREE_BUILD_TESTS) - return() - endif() - - cmake_parse_arguments( - _RULE - "" - "NAME" - "SRCS;DEPS;LABELS" - ${ARGN} - ) - - iree_package_name(_PACKAGE_NAME) - set(_NAME "${_PACKAGE_NAME}_${_RULE_NAME}") - - iree_package_ns(_PACKAGE_NS) - string(REPLACE "::" "/" _PACKAGE_PATH ${_PACKAGE_NS}) - set(_NAME_PATH "${_PACKAGE_PATH}:${_RULE_NAME}") - - add_test( - NAME ${_NAME_PATH} - COMMAND ${CMAKE_SOURCE_DIR}/build_tools/cmake/run_test.sh ${PYTHON_EXECUTABLE} "${CMAKE_CURRENT_SOURCE_DIR}/${_RULE_SRCS}" - WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR} - ) - - list(APPEND _RULE_LABELS "${_PACKAGE_PATH}") - set_property(TEST ${_NAME_PATH} PROPERTY LABELS "${_RULE_LABELS}") - set_property(TEST ${_NAME_PATH} PROPERTY ENVIRONMENT "PYTHONPATH=${CMAKE_BINARY_DIR}/bindings/python:$ENV{PYTHONPATH};TEST_TMPDIR=${_NAME}_test_tmpdir") - # TODO(marbre): Find out how to add deps to tests. - # Similar to _RULE_DATA in iree_lit_test(). - -endfunction()
diff --git a/bindings/python/build_tools/cmake/iree_pybind_cc_library.cmake b/bindings/python/build_tools/cmake/iree_pybind_cc_library.cmake deleted file mode 100644 index 589f768..0000000 --- a/bindings/python/build_tools/cmake/iree_pybind_cc_library.cmake +++ /dev/null
@@ -1,127 +0,0 @@ -# Copyright 2020 Google LLC -# -# Licensed under the Apache License, Version 2.0 (the "License"); -# you may not use this file except in compliance with the License. -# You may obtain a copy of the License at -# -# https://www.apache.org/licenses/LICENSE-2.0 -# -# Unless required by applicable law or agreed to in writing, software -# distributed under the License is distributed on an "AS IS" BASIS, -# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -# See the License for the specific language governing permissions and -# limitations under the License. - -include(CMakeParseArguments) - -# iree_pybind_cc_library() -# -# CMake function to imitate Bazel's pybind_cc_library rule. -# -# Parameters: -# NAME: name of target -# HDRS: List of public header files for the library -# SRCS: List of source files for the library -# DEPS: List of other libraries to be linked in to the binary targets -# COPTS: List of private compile options -# DEFINES: List of public defines -# INCLUDES: Include directories to add to dependencies -# LINKOPTS: List of link options -# TYPE: Type of library to be crated: either "MODULE", "SHARED" or "STATIC" (default). -# TESTONLY: When added, this target will only be built if user passes -DIREE_BUILD_TESTS=ON to CMake. - -function(iree_pybind_cc_library) - - cmake_parse_arguments( - _RULE - "TESTONLY" - "NAME" - "HDRS;SRCS;COPTS;DEFINES;LINKOPTS;DEPS;INCLUDES;TYPE" - ${ARGN} - ) - - iree_package_ns(_PACKAGE_NS) - # Replace dependencies passed by ::name with ::iree::package::name - list(TRANSFORM _RULE_DEPS REPLACE "^::" "${_PACKAGE_NS}::") - - if(NOT _RULE_TESTONLY OR IREE_BUILD_TESTS) - # Prefix the library with the package name, so we get: iree_package_name. - iree_package_name(_PACKAGE_NAME) - set(_NAME "${_PACKAGE_NAME}_${_RULE_NAME}") - - if(NOT _RULE_TYPE) - set(_RULE_TYPE "STATIC") - endif() - - string(TOUPPER "${_RULE_TYPE}" _uppercase_RULE_TYPE) - if(NOT _uppercase_RULE_TYPE MATCHES "^(STATIC|SHARED|MODULE)") - message(FATAL_ERROR "Unsported library TYPE for iree_pybind_cc_library: ${_RULE_TYPE}") - endif() - - add_library(${_NAME} ${_uppercase_RULE_TYPE} "") - target_sources(${_NAME} - PRIVATE - ${_RULE_SRCS} - ${_RULE_HDRS} - ) - target_include_directories(${_NAME} - PUBLIC - "$<BUILD_INTERFACE:${IREE_COMMON_INCLUDE_DIRS}>" - "$<BUILD_INTERFACE:${_RULE_INCLUDES}>" - PRIVATE - ${PYBIND11_INCLUDE_DIR} - ${PYTHON_INCLUDE_DIR} - ) - target_compile_options(${_NAME} - PRIVATE - ${_RULE_COPTS} - ${PYBIND_COPTS} - ${IREE_DEFAULT_COPTS} - ) - - target_link_libraries(${_NAME} - PUBLIC - ${_RULE_DEPS} - PRIVATE - pybind11 - ${PYTHON_LIBRARY} - ${_RULE_LINKOPTS} - ${IREE_DEFAULT_LINKOPTS} - ) - target_compile_definitions(${_NAME} - PUBLIC - ${_RULE_DEFINES} - ) - - # Add all IREE targets to a folder in the IDE for organization. - if(_RULE_PUBLIC) - set_property(TARGET ${_NAME} PROPERTY FOLDER ${IREE_IDE_FOLDER}) - elseif(_RULE_TESTONLY) - set_property(TARGET ${_NAME} PROPERTY FOLDER ${IREE_IDE_FOLDER}/test) - else() - set_property(TARGET ${_NAME} PROPERTY FOLDER ${IREE_IDE_FOLDER}/internal) - endif() - - # INTERFACE libraries can't have the CXX_STANDARD property set. - set_property(TARGET ${_NAME} PROPERTY CXX_STANDARD ${IREE_CXX_STANDARD}) - set_property(TARGET ${_NAME} PROPERTY CXX_STANDARD_REQUIRED ON) - - set_target_properties(${_NAME} PROPERTIES OUTPUT_NAME "${_RULE_NAME}") - - if(NOT _uppercase_RULE_TYPE MATCHES "STATIC") - set_property(TARGET ${_NAME} PROPERTY PREFIX "${PYTHON_MODULE_PREFIX}") - set_property(TARGET ${_NAME} PROPERTY SUFFIX "${PYTHON_MODULE_EXTENSION}") - endif() - - # Alias the iree_package_name library to iree::package::name. - # This lets us more clearly map to Bazel and makes it possible to - # disambiguate the underscores in paths vs. the separators. - add_library(${_PACKAGE_NS}::${_RULE_NAME} ALIAS ${_NAME}) - iree_package_dir(_PACKAGE_DIR) - if(${_RULE_NAME} STREQUAL ${_PACKAGE_DIR}) - # If the library name matches the package then treat it as a default. - # For example, foo/bar/ library 'bar' would end up as 'foo::bar'. - add_library(${_PACKAGE_NS} ALIAS ${_NAME}) - endif() - endif() -endfunction()
diff --git a/bindings/python/pyiree/common/.skip_bazel_to_cmake b/bindings/python/pyiree/common/.skip_bazel_to_cmake new file mode 100644 index 0000000..e69de29 --- /dev/null +++ b/bindings/python/pyiree/common/.skip_bazel_to_cmake
diff --git a/bindings/python/pyiree/common/CMakeLists.txt b/bindings/python/pyiree/common/CMakeLists.txt index 19f318a..a067903 100644 --- a/bindings/python/pyiree/common/CMakeLists.txt +++ b/bindings/python/pyiree/common/CMakeLists.txt
@@ -12,13 +12,12 @@ # See the License for the specific language governing permissions and # limitations under the License. -iree_pybind_cc_library( +iree_pyext_library( NAME - common - HDRS + PyextCommonLib + SRCS "binding.h" "status_utils.h" - SRCS "status_utils.cc" DEPS iree::base::api
diff --git a/bindings/python/pyiree/compiler/.skip_bazel_to_cmake b/bindings/python/pyiree/compiler/.skip_bazel_to_cmake new file mode 100644 index 0000000..e69de29 --- /dev/null +++ b/bindings/python/pyiree/compiler/.skip_bazel_to_cmake
diff --git a/bindings/python/pyiree/compiler/CMakeLists.txt b/bindings/python/pyiree/compiler/CMakeLists.txt index 35dbc6f..b7bc2cc 100644 --- a/bindings/python/pyiree/compiler/CMakeLists.txt +++ b/bindings/python/pyiree/compiler/CMakeLists.txt
@@ -17,41 +17,41 @@ compiler SRCS "__init__.py" - DEPS - ::binding + PYEXT_DEPS + ::PyExtCompiler ) -iree_py_extension( +iree_pyext_module( NAME + PyExtCompiler + MODULE_NAME binding SRCS "initialize_module.cc" - DEPS - ::compiler_library - bindings::python::pyiree::common - COPTS - ${PYBIND_COPTS} - ${PYBIND_EXTENSION_COPTS} - TYPE - SHARED + PYEXT_DEPS + ::PyExtCompilerLib + bindings::python::pyiree::common::PyextCommonLib ) -iree_pybind_cc_library( +iree_pyext_library( NAME - compiler_library - HDRS - "compiler.h" + PyExtCompilerLib SRCS + "compiler.h" "compiler.cc" + PYEXT_DEPS + bindings::python::pyiree::common::PyextCommonLib + COPTS + ${PYBIND_REGISTER_MLIR_PASSES} DEPS - # Transforms. Adopted from the Bazel variable COMPILER_DEPS. + # Transforms. Adopted from the Bazel variable COMPILER::DEPS. iree::compiler::Dialect::Flow::Transforms iree::compiler::Dialect::HAL::Transforms iree::compiler::Dialect::HAL::Target iree::compiler::Dialect::Shape::IR iree::compiler::Dialect::Shape::Transforms iree::compiler::Dialect::VM::Transforms - # Targets. Adopted from the Bazel variable COMPILER_DEPS. + # Targets. Adopted from the Bazel variable COMPILER::DEPS. iree::compiler::Dialect::HAL::Target::VMLA iree::compiler::Dialect::HAL::Target::LLVM::LLVMAOT iree::compiler::Dialect::HAL::Target::LLVM::LLVMIR @@ -62,16 +62,11 @@ iree::tools::init_iree_passes_and_dialects iree::tools::init_mlir_passes_and_dialects iree::tools::init_targets - bindings::python::pyiree::common LLVMSupport MLIRIR MLIRSCFTransforms MLIRParser MLIRPass - COPTS - ${PYBIND_REGISTER_MLIR_PASSES} - TYPE - STATIC ) iree_py_test(
diff --git a/bindings/python/pyiree/rt/.skip_bazel_to_cmake b/bindings/python/pyiree/rt/.skip_bazel_to_cmake new file mode 100644 index 0000000..e69de29 --- /dev/null +++ b/bindings/python/pyiree/rt/.skip_bazel_to_cmake
diff --git a/bindings/python/pyiree/rt/BUILD b/bindings/python/pyiree/rt/BUILD index ac5a93b..6d77a7f 100644 --- a/bindings/python/pyiree/rt/BUILD +++ b/bindings/python/pyiree/rt/BUILD
@@ -92,9 +92,9 @@ "//iree/vm", "//iree/vm:bytecode_module", "//iree/vm:invocation", + "//iree/vm:list", "//iree/vm:module", "//iree/vm:ref", - "//iree/vm:variant_list", "@com_google_absl//absl/container:inlined_vector", "@com_google_absl//absl/memory", "@com_google_absl//absl/strings",
diff --git a/bindings/python/pyiree/rt/CMakeLists.txt b/bindings/python/pyiree/rt/CMakeLists.txt index cf883ac..1ec5eaf 100644 --- a/bindings/python/pyiree/rt/CMakeLists.txt +++ b/bindings/python/pyiree/rt/CMakeLists.txt
@@ -12,50 +12,21 @@ # See the License for the specific language governing permissions and # limitations under the License. -iree_py_library( +iree_pyext_library( NAME - rt + PyExtRtLib SRCS - "__init__.py" - # "system_api.py" - DEPS - ::binding -) - -iree_py_extension( - NAME - binding - SRCS - "initialize_module.cc" - DEPS - iree::hal::vulkan::vulkan_driver_module - iree::hal::llvmjit::llvmjit_driver_module - iree::hal::vmla::vmla_driver_module - ::rt_library - bindings::python::pyiree::common - iree::base::initializer - COPTS - ${PYBIND_COPTS} - ${PYBIND_EXTENSION_COPTS} - TYPE - SHARED -) - -iree_pybind_cc_library( - NAME - rt_library - HDRS "function_abi.h" "hal.h" "host_types.h" "vm.h" - SRCS "function_abi.cc" "hal.cc" "host_types.cc" "vm.cc" + PYEXT_DEPS + bindings::python::pyiree::common::PyextCommonLib DEPS - bindings::python::pyiree::common iree::base::api iree::base::signature_mangle iree::hal::api @@ -66,23 +37,37 @@ iree::vm::bytecode_module iree::vm::invocation iree::vm::ref - iree::vm::variant_list absl::inlined_vector absl::memory absl::strings absl::optional absl::span - TYPE - STATIC +) + +iree_pyext_module( + NAME + PyExtRt + MODULE_NAME binding + SRCS + "initialize_module.cc" + PYEXT_DEPS + ::PyExtRtLib + bindings::python::pyiree::common::PyextCommonLib + DEPS + iree::hal::vulkan::vulkan_driver_module + iree::hal::llvmjit::llvmjit_driver_module + iree::hal::vmla::vmla_driver_module + iree::base::initializer ) iree_py_library( NAME - system_api + rt SRCS + "__init__.py" "system_api.py" - DEPS - ::binding + PYEXT_DEPS + ::PyExtRt ) iree_py_test(
diff --git a/bindings/python/pyiree/rt/function_abi.cc b/bindings/python/pyiree/rt/function_abi.cc index 6b8c01d..58a0295 100644 --- a/bindings/python/pyiree/rt/function_abi.cc +++ b/bindings/python/pyiree/rt/function_abi.cc
@@ -24,8 +24,8 @@ #include "iree/base/signature_mangle.h" #include "iree/hal/api.h" #include "iree/modules/hal/hal_module.h" +#include "iree/vm/list.h" #include "iree/vm/ref.h" -#include "iree/vm/variant_list.h" namespace iree { namespace python { @@ -184,12 +184,12 @@ default: throw RaisePyError(PyExc_NotImplementedError, "Unsupported scalar type"); } - CheckApiStatus(iree_vm_variant_list_append_value(f_args.raw_ptr(), value), + CheckApiStatus(iree_vm_list_push_value(f_args.raw_ptr(), &value), "Could not pack scalar argument"); } py::object UnpackScalar(const RawSignatureParser::Description& desc, - iree_vm_variant_t& f_result) { + const iree_vm_variant_t& f_result) { switch (desc.scalar.type) { case AbiConstants::ScalarType::kUint8: case AbiConstants::ScalarType::kUint16: @@ -296,12 +296,15 @@ } for (size_t i = 0, e = descs.size(); i < e; ++i) { const Description& desc = descs[i]; - iree_vm_variant_t* f_result = - iree_vm_variant_list_get(f_results.raw_ptr(), i); + iree_vm_variant_t f_result = iree_vm_variant_empty(); + if (!iree_status_is_ok( + iree_vm_list_get_variant(f_results.raw_ptr(), i, &f_result))) { + throw RaiseValueError("Could not get result from list"); + } switch (desc.type) { case RawSignatureParser::Type::kBuffer: { iree_hal_buffer_view_t* buffer_view = - iree_hal_buffer_view_deref(&f_result->ref); + iree_hal_buffer_view_deref(&f_result.ref); if (!buffer_view) { throw RaiseValueError( "Could not deref result buffer view (wrong type?)"); @@ -340,7 +343,7 @@ "Ref objects not yet supported"); break; case RawSignatureParser::Type::kScalar: - py_results[i] = UnpackScalar(desc, *f_result); + py_results[i] = UnpackScalar(desc, f_result); break; default: throw RaisePyError(PyExc_NotImplementedError, @@ -397,9 +400,9 @@ iree_hal_buffer_release(raw_buffer); iree_vm_ref_t buffer_view_ref = iree_hal_buffer_view_move_ref(buffer_view); - CheckApiStatus(iree_vm_variant_list_append_ref_move(f_results.raw_ptr(), - &buffer_view_ref), - "Error moving buffer"); + CheckApiStatus( + iree_vm_list_push_ref_move(f_results.raw_ptr(), &buffer_view_ref), + "Error moving buffer"); break; } case RawSignatureParser::Type::kRefObject: @@ -484,9 +487,8 @@ "Error allocating buffer_view"); iree_hal_buffer_release(raw_buffer); iree_vm_ref_t buffer_view_ref = iree_hal_buffer_view_move_ref(buffer_view); - CheckApiStatus( - iree_vm_variant_list_append_ref_move(f_args.raw_ptr(), &buffer_view_ref), - "Error moving buffer view"); + CheckApiStatus(iree_vm_list_push_ref_move(f_args.raw_ptr(), &buffer_view_ref), + "Error moving buffer view"); } void SetupFunctionAbiBindings(pybind11::module m) {
diff --git a/bindings/python/pyiree/rt/function_abi_test.py b/bindings/python/pyiree/rt/function_abi_test.py index 9f34fb6..cb8c804 100644 --- a/bindings/python/pyiree/rt/function_abi_test.py +++ b/bindings/python/pyiree/rt/function_abi_test.py
@@ -49,9 +49,15 @@ def setUpClass(cls): super().setUpClass() driver_names = rt.HalDriver.query() - print("DRIVER_NAMES =", driver_names) - cls.driver = rt.HalDriver.create("vulkan") - cls.device = cls.driver.create_default_device() + for driver_name in driver_names: + print("Try create driver:", driver_name) + try: + cls.driver = rt.HalDriver.create(driver_name) + cls.device = cls.driver.create_default_device() + except Exception: + print("Could not create driver:", driver_name) + else: + break def setUp(self): super().setUp()
diff --git a/bindings/python/pyiree/rt/system_api.py b/bindings/python/pyiree/rt/system_api.py index b7deef6..aaea01f 100644 --- a/bindings/python/pyiree/rt/system_api.py +++ b/bindings/python/pyiree/rt/system_api.py
@@ -64,9 +64,6 @@ continue try: driver = _binding.HalDriver.create(driver_name) - print( - "Created IREE driver %s: %r" % (driver_name, driver), file=sys.stderr) - return driver # TODO(laurenzo): Remove these prints to stderr (for now, more information # is better and there is no better way to report it yet). except Exception as ex: # pylint: disable=broad-except @@ -74,6 +71,23 @@ "Could not create default driver %s: %r" % (driver_name, ex), file=sys.stderr) driver_exceptions[driver_name] = ex + continue + + # Sanity check creation of the default device and skip the driver if + # this fails (this works around issues where the driver is present + # but there are no devices). This default initialization scheme needs + # to be improved. + try: + device = driver.create_default_device() + except Exception as ex: + print( + "Could not create default driver device %s: %r" % (driver_name, ex), + file=sys.stderr) + driver_exceptions[driver_name] = ex + continue + + print("Created IREE driver %s: %r" % (driver_name, driver), file=sys.stderr) + return driver # All failed. raise RuntimeError("Could not create any requested driver "
diff --git a/bindings/python/pyiree/rt/vm.cc b/bindings/python/pyiree/rt/vm.cc index 9674b16..0f79099 100644 --- a/bindings/python/pyiree/rt/vm.cc +++ b/bindings/python/pyiree/rt/vm.cc
@@ -197,21 +197,25 @@ absl::StrAppend(&s, "<VmVariantList(", size(), "): ["); for (iree_host_size_t i = 0, e = size(); i < e; ++i) { - iree_vm_variant_t* variant = - iree_vm_variant_list_get(mutable_this->raw_ptr(), i); + iree_vm_variant_t variant = iree_vm_variant_empty(); + if (!iree_status_is_ok( + iree_vm_list_get_variant(mutable_this->raw_ptr(), i, &variant))) { + absl::StrAppend(&s, "Error"); + continue; + } if (i > 0) absl::StrAppend(&s, ", "); - if (IREE_VM_VARIANT_IS_VALUE(variant)) { - absl::StrAppend(&s, variant->i32); - } else if (IREE_VM_VARIANT_IS_REF(variant)) { + if (iree_vm_variant_is_value(variant)) { + absl::StrAppend(&s, variant.i32); + } else if (iree_vm_variant_is_ref(variant)) { // Pretty print a subset of ABI impacting known types. - if (iree_hal_buffer_isa(&variant->ref)) { - auto* hal_buffer = iree_hal_buffer_deref(&variant->ref); + if (iree_hal_buffer_isa(&variant.ref)) { + auto* hal_buffer = iree_hal_buffer_deref(&variant.ref); assert(hal_buffer); absl::StrAppend(&s, "HalBuffer(", iree_hal_buffer_byte_length(hal_buffer), ")"); - } else if (iree_hal_buffer_view_isa(&variant->ref)) { - auto hal_bv = iree_hal_buffer_view_deref(&variant->ref); + } else if (iree_hal_buffer_view_isa(&variant.ref)) { + auto hal_bv = iree_hal_buffer_view_deref(&variant.ref); absl::StrAppend(&s, "HalBufferView("); absl::InlinedVector<int32_t, 5> shape( iree_hal_buffer_view_shape_rank(hal_bv)); @@ -221,7 +225,7 @@ iree_hal_buffer_view_element_type(hal_bv))), ")"); } else { - absl::StrAppend(&s, "Unknown(", variant->ref_type, ")"); + absl::StrAppend(&s, "Unknown(", variant.type.ref_type, ")"); } } else { absl::StrAppend(&s, "None");
diff --git a/bindings/python/pyiree/rt/vm.h b/bindings/python/pyiree/rt/vm.h index e281f2a..1e7f984 100644 --- a/bindings/python/pyiree/rt/vm.h +++ b/bindings/python/pyiree/rt/vm.h
@@ -21,7 +21,7 @@ #include "iree/base/api.h" #include "iree/vm/api.h" #include "iree/vm/bytecode_module.h" -#include "iree/vm/variant_list.h" +#include "iree/vm/list.h" namespace iree { namespace python { @@ -67,7 +67,7 @@ VmVariantList() : list_(nullptr) {} ~VmVariantList() { if (list_) { - iree_vm_variant_list_free(list_); + iree_vm_list_release(list_); } } @@ -80,28 +80,29 @@ VmVariantList(const VmVariantList&) = delete; static VmVariantList Create(iree_host_size_t capacity) { - iree_vm_variant_list_t* list; - CheckApiStatus( - iree_vm_variant_list_alloc(capacity, IREE_ALLOCATOR_SYSTEM, &list), - "Error allocating variant list"); + iree_vm_list_t* list; + CheckApiStatus(iree_vm_list_create(/*element_type=*/nullptr, capacity, + IREE_ALLOCATOR_SYSTEM, &list), + "Error allocating variant list"); return VmVariantList(list); } - iree_host_size_t size() const { return iree_vm_variant_list_size(list_); } + iree_host_size_t size() const { return iree_vm_list_size(list_); } - iree_vm_variant_list_t* raw_ptr() { return list_; } - const iree_vm_variant_list_t* raw_ptr() const { return list_; } + iree_vm_list_t* raw_ptr() { return list_; } + const iree_vm_list_t* raw_ptr() const { return list_; } void AppendNullRef() { - CheckApiStatus(iree_vm_variant_list_append_null_ref(raw_ptr()), + iree_vm_ref_t null_ref = {0}; + CheckApiStatus(iree_vm_list_push_ref_move(raw_ptr(), &null_ref), "Error appending to list"); } std::string DebugString() const; private: - VmVariantList(iree_vm_variant_list_t* list) : list_(list) {} - iree_vm_variant_list_t* list_; + VmVariantList(iree_vm_list_t* list) : list_(list) {} + iree_vm_list_t* list_; }; //------------------------------------------------------------------------------
diff --git a/build_tools/bazel/iree_flatcc.bzl b/build_tools/bazel/iree_flatcc.bzl new file mode 100644 index 0000000..e091af6 --- /dev/null +++ b/build_tools/bazel/iree_flatcc.bzl
@@ -0,0 +1,57 @@ +# Copyright 2020 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +"""Generates flatbuffer source files with flatcc.""" + +def iree_flatbuffer_c_library( + name, + srcs, + flatcc_args = ["--common", "--reader"], + testonly = False, + **kwargs): + flatcc = "@com_github_dvidelabs_flatcc//:flatcc" + flatcc_rt = "@com_github_dvidelabs_flatcc//:runtime" + + flags = [ + "-o$(RULEDIR)", + ] + flatcc_args + + out_stem = "%s" % (srcs[0].replace(".fbs", "")) + + outs = [] + for arg in flags: + if arg == "--reader": + outs += ["%s_reader.h" % (out_stem)] + if arg == "--builder": + outs += ["%s_builder.h" % (out_stem)] + if arg == "--verifier": + outs += ["%s_verifier.h" % (out_stem)] + + native.genrule( + name = name + "_gen", + srcs = srcs, + outs = outs, + tools = [flatcc], + cmd = "$(location %s) %s $(SRCS)" % (flatcc, " ".join(flags)), + testonly = testonly, + ) + native.cc_library( + name = name, + hdrs = outs, + deps = [ + flatcc_rt, + ], + testonly = testonly, + **kwargs + )
diff --git a/build_tools/bazel/third_party_import/llvm-project/overlay/llvm/BUILD.bazel b/build_tools/bazel/third_party_import/llvm-project/overlay/llvm/BUILD.bazel index bade7ab..befc20c 100644 --- a/build_tools/bazel/third_party_import/llvm-project/overlay/llvm/BUILD.bazel +++ b/build_tools/bazel/third_party_import/llvm-project/overlay/llvm/BUILD.bazel
@@ -155,10 +155,10 @@ name = "InstCombineTableGen", tbl_outs = [( "-gen-searchable-tables", - "lib/Transforms/InstCombine/InstCombineTables.inc", + "lib/Target/AMDGPU/InstCombineTables.inc", )], tblgen = ":llvm-tblgen", - td_file = "lib/Transforms/InstCombine/InstCombineTables.td", + td_file = "lib/Target/AMDGPU/InstCombineTables.td", td_srcs = glob([ "include/llvm/CodeGen/*.td", "include/llvm/IR/Intrinsics*.td", @@ -721,6 +721,7 @@ "lib/Analysis/*.h", ], exclude = [ + "lib/Analysis/DevelopmentModeInlineAdvisor.cpp", "lib/Analysis/MLInlineAdvisor.cpp", "lib/Analysis/ReleaseModeModelRunner.cpp", "lib/Analysis/TFUtils.cpp", @@ -3187,6 +3188,7 @@ ]), copts = llvm_copts, deps = [ + ":BinaryFormat", ":DebugInfoCodeView", ":MC", ":Object",
diff --git a/build_tools/bazel/third_party_import/llvm-project/overlay/mlir/BUILD.bazel b/build_tools/bazel/third_party_import/llvm-project/overlay/mlir/BUILD.bazel index ec0574f..ae413a1 100644 --- a/build_tools/bazel/third_party_import/llvm-project/overlay/mlir/BUILD.bazel +++ b/build_tools/bazel/third_party_import/llvm-project/overlay/mlir/BUILD.bazel
@@ -387,7 +387,7 @@ "include/mlir/Interfaces/CallInterfaces.td", "include/mlir/Interfaces/ControlFlowInterfaces.td", "include/mlir/Interfaces/SideEffectInterfaces.td", - "include/mlir/Interfaces/VectorUnrollInterface.td", + "include/mlir/Interfaces/VectorInterfaces.td", "include/mlir/Interfaces/ViewLikeInterface.td", ":OpBaseTdFiles", ], @@ -500,6 +500,7 @@ deps = [ ":Affine", ":IR", + ":Support", "@llvm-project//llvm:Support", ], ) @@ -647,13 +648,13 @@ ) cc_library( - name = "VectorUnrollInterface", - srcs = ["lib/Interfaces/VectorUnrollInterface.cpp"], - hdrs = ["include/mlir/Interfaces/VectorUnrollInterface.h"], + name = "VectorInterfaces", + srcs = ["lib/Interfaces/VectorInterfaces.cpp"], + hdrs = ["include/mlir/Interfaces/VectorInterfaces.h"], includes = ["include"], deps = [ ":IR", - ":VectorUnrollInterfaceIncGen", + ":VectorInterfacesIncGen", ], ) @@ -855,7 +856,7 @@ ":SideEffectInterfaces", ":StandardOpsIncGen", ":Support", - ":VectorUnrollInterface", + ":VectorInterfaces", ":ViewLikeInterface", "@llvm-project//llvm:Support", ], @@ -918,9 +919,9 @@ ":SideEffectInterfaces", ":StandardOps", ":Support", + ":VectorInterfaces", ":VectorOpsIncGen", ":VectorTransformPatternsIncGen", - ":VectorUnrollInterface", "@llvm-project//llvm:Support", ], ) @@ -2127,20 +2128,20 @@ ) gentbl( - name = "VectorUnrollInterfaceIncGen", + name = "VectorInterfacesIncGen", strip_include_prefix = "include", tbl_outs = [ ( "-gen-op-interface-decls", - "include/mlir/Interfaces/VectorUnrollInterface.h.inc", + "include/mlir/Interfaces/VectorInterfaces.h.inc", ), ( "-gen-op-interface-defs", - "include/mlir/Interfaces/VectorUnrollInterface.cpp.inc", + "include/mlir/Interfaces/VectorInterfaces.cpp.inc", ), ], tblgen = ":mlir-tblgen", - td_file = "include/mlir/Interfaces/VectorUnrollInterface.td", + td_file = "include/mlir/Interfaces/VectorInterfaces.td", td_srcs = [ ":OpBaseTdFiles", ], @@ -3271,6 +3272,7 @@ ":QuantPassIncGen", ":SideEffectInterfaces", ":StandardOps", + ":TransformUtils", "@llvm-project//llvm:Support", ], ) @@ -3586,7 +3588,7 @@ name = "VectorOpsTdFiles", srcs = [ "include/mlir/Dialect/Vector/VectorOps.td", - "include/mlir/Interfaces/VectorUnrollInterface.td", + "include/mlir/Interfaces/VectorInterfaces.td", ":AffineOpsTdFiles", ":OpBaseTdFiles", ], @@ -3716,7 +3718,7 @@ "include/mlir/Interfaces/ControlFlowInterfaces.h", "include/mlir/Interfaces/ControlFlowInterfaces.td", "include/mlir/Interfaces/SideEffectInterfaces.td", - "include/mlir/Interfaces/VectorUnrollInterface.td", + "include/mlir/Interfaces/VectorInterfaces.td", "include/mlir/Interfaces/ViewLikeInterface.td", "include/mlir/Dialect/LLVMIR/LLVMOpBase.td", "include/mlir/Dialect/StandardOps/IR/Ops.td",
diff --git a/build_tools/bazel_to_cmake/bazel_to_cmake.py b/build_tools/bazel_to_cmake/bazel_to_cmake.py index ed70b9c..51e559c 100755 --- a/build_tools/bazel_to_cmake/bazel_to_cmake.py +++ b/build_tools/bazel_to_cmake/bazel_to_cmake.py
@@ -103,11 +103,12 @@ if not os.path.isdir(directory_path): raise FileNotFoundError(f"Cannot find directory '{directory_path}'") + skip_file_path = os.path.join(directory_path, ".skip_bazel_to_cmake") build_file_path = os.path.join(directory_path, "BUILD") cmakelists_file_path = os.path.join(directory_path, "CMakeLists.txt") - if not os.path.isfile(build_file_path): - # No Bazel BUILD file in this directory to convert, skip. + if os.path.isfile(skip_file_path) or not os.path.isfile(build_file_path): + # No Bazel BUILD file in this directory or explicit skip. return global repo_root
diff --git a/build_tools/bazel_to_cmake/bazel_to_cmake_converter.py b/build_tools/bazel_to_cmake/bazel_to_cmake_converter.py index 34e09a8..bc40c97 100644 --- a/build_tools/bazel_to_cmake/bazel_to_cmake_converter.py +++ b/build_tools/bazel_to_cmake/bazel_to_cmake_converter.py
@@ -234,6 +234,13 @@ flatc_args = "\n".join([f' "{flatc_arg}"' for flatc_arg in flatc_args]) return f" FLATC_ARGS\n{flatc_args}\n" + def _convert_flatcc_args_block(self, flatcc_args): + if not flatcc_args: + return "" + flatcc_args = "\n".join( + [f' "{flatcc_arg}"' for flatcc_arg in flatcc_args]) + return f" FLATCC_ARGS\n{flatcc_args}\n" + def _convert_unimplemented_function(self, function, details=""): message = f"Unimplemented {function}: {details}" if not self.converter.first_error: @@ -481,6 +488,17 @@ f"{flags_block}" f" PUBLIC\n)\n\n") + def iree_flatbuffer_c_library(self, name, srcs, flatcc_args=None): + name_block = self._convert_name_block(name) + srcs_block = self._convert_srcs_block(srcs) + flatcc_args_block = self._convert_flatcc_args_block(flatcc_args) + + self.converter.body += (f"flatbuffer_c_library(\n" + f"{name_block}" + f"{srcs_block}" + f"{flatcc_args_block}" + f" PUBLIC\n)\n\n") + def iree_flatbuffer_cc_library(self, name, srcs, flatc_args=None): name_block = self._convert_name_block(name) srcs_block = self._convert_srcs_block(srcs)
diff --git a/build_tools/bazel_to_cmake/bazel_to_cmake_targets.py b/build_tools/bazel_to_cmake/bazel_to_cmake_targets.py index 7f9b497..19a0fd3 100644 --- a/build_tools/bazel_to_cmake/bazel_to_cmake_targets.py +++ b/build_tools/bazel_to_cmake/bazel_to_cmake_targets.py
@@ -64,6 +64,8 @@ # Misc single targets "@com_google_benchmark//:benchmark": ["benchmark"], "@com_github_google_flatbuffers//:flatbuffers": ["flatbuffers"], + "@com_github_dvidelabs_flatcc//:flatcc": ["flatcc"], + "@com_github_dvidelabs_flatcc//:runtime": ["flatcc::runtime"], "@com_google_googletest//:gtest": ["gmock", "gtest"], "@renderdoc_api//:renderdoc_app": ["renderdoc_api::renderdoc_app"], "@sdl2//:SDL2": ["SDL2-static"]
diff --git a/build_tools/cmake/build_docs.sh b/build_tools/cmake/build_docs.sh index d4a3c84..191ed17 100755 --- a/build_tools/cmake/build_docs.sh +++ b/build_tools/cmake/build_docs.sh
@@ -49,34 +49,12 @@ cd ${ROOT_DIR?} +cp README.md ${BUILD_DIR}/doc/index.md +cp -rf docs/* ${BUILD_DIR}/doc/ + # Update op_coverage.md scripts/update_op_coverage.py ${BUILD_DIR} # Update e2e_coverage.md PYTHON_BIN=`which python3` scripts/update_e2e_coverage.py ${BUILD_DIR} -# Copy a curated list of docs to publish. This is expected to cover all docs -# under docs/ after they are refreshed. - -cp README.md ${BUILD_DIR}/doc/index.md -cp docs/IREE-Architecture.svg ${BUILD_DIR}/doc/ - -cp docs/roadmap.md ${BUILD_DIR}/doc/ -cp docs/roadmap_design.md ${BUILD_DIR}/doc/ -cp docs/developer_overview.md ${BUILD_DIR}/doc/ -cp docs/testing_guide.md ${BUILD_DIR}/doc/ -cp docs/iree_community.md ${BUILD_DIR}/doc/ - -mkdir -p ${BUILD_DIR}/doc/GetStarted/ -cp docs/GetStarted/getting_started_windows_bazel.md ${BUILD_DIR}/doc/GetStarted/ -cp docs/GetStarted/getting_started_windows_cmake.md ${BUILD_DIR}/doc/GetStarted/ -cp docs/GetStarted/getting_started_windows_vulkan.md ${BUILD_DIR}/doc/GetStarted/ -cp docs/GetStarted/getting_started_linux_bazel.md ${BUILD_DIR}/doc/GetStarted/ -cp docs/GetStarted/getting_started_linux_cmake.md ${BUILD_DIR}/doc/GetStarted/ -cp docs/GetStarted/getting_started_linux_vulkan.md ${BUILD_DIR}/doc/GetStarted/ -cp docs/GetStarted/getting_started_macos_bazel.md ${BUILD_DIR}/doc/GetStarted/ -cp docs/GetStarted/getting_started_macos_cmake.md ${BUILD_DIR}/doc/GetStarted/ -cp docs/GetStarted/getting_started_android_cmake.md ${BUILD_DIR}/doc/GetStarted/ -cp docs/GetStarted/getting_started_python.md ${BUILD_DIR}/doc/GetStarted/ -cp docs/GetStarted/generic_vulkan_env_setup.md ${BUILD_DIR}/doc/GetStarted/ -cp docs/GetStarted/cmake_options_and_variables.md ${BUILD_DIR}/doc/GetStarted/
diff --git a/build_tools/cmake/flatbuffer_c_library.cmake b/build_tools/cmake/flatbuffer_c_library.cmake new file mode 100644 index 0000000..916cfd9 --- /dev/null +++ b/build_tools/cmake/flatbuffer_c_library.cmake
@@ -0,0 +1,158 @@ +# Copyright 2020 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +include(CMakeParseArguments) + +# flatbuffer_c_library() +# +# CMake function to invoke the flatcc compiler. +# +# Parameters: +# NAME: name of target (see Note) +# SRCS: List of source files for the library +# DEPS: List of other libraries to be linked in to the binary targets +# COPTS: List of private compile options +# DEFINES: List of public defines +# LINKOPTS: List of link options +# FLATCC_ARGS: List of flattbuffers arguments. Default: +# "--common" +# "--reader" +# PUBLIC: Add this so that this library will be exported under iree:: +# Also in IDE, target will appear in IREE folder while non PUBLIC will be in IREE/internal. +# TESTONLY: When added, this target will only be built if user passes -DIREE_BUILD_TESTS=ON to CMake. +# +# Note: +# By default, flatbuffer_c_library will always create a library named ${NAME}, +# and alias target iree::${NAME}. The iree:: form should always be used. +# This is to reduce namespace pollution. +# +# flatbuffer_c_library( +# NAME +# base_schema +# SRCS +# "a.fbs" +# ) +# flatbuffer_c_library( +# NAME +# other_schemas +# SRCS +# "b.fbs" +# DEPS +# iree::schemas::base_schema +# PUBLIC +# ) +# +# iree_cc_binary( +# NAME +# main_lib +# ... +# DEPS +# iree::schemas::other_schemas +# ) +function(flatbuffer_c_library) + cmake_parse_arguments(_RULE + "PUBLIC;TESTONLY" + "NAME" + "SRCS;COPTS;DEFINES;LINKOPTS;DEPS;FLATCC_ARGS" + ${ARGN} + ) + + if(_RULE_TESTONLY AND NOT IREE_BUILD_TESTS) + return() + endif() + + # Prefix the library with the package name, so we get: iree_package_name + iree_package_name(_PACKAGE_NAME) + set(_NAME "${_PACKAGE_NAME}_${_RULE_NAME}") + + if(NOT DEFINED _RULE_FLATCC_ARGS) + set(_RULE_FLATCC_ARGS + "--common" + "--reader" + ) + else() + set(_RULE_FLATCC_ARGS ${_RULE_FLATCC_ARGS}) + endif() + + set(_OUTS "") + foreach(_SRC ${_RULE_SRCS}) + get_filename_component(_SRC_FILENAME ${_SRC} NAME_WE) + foreach(_ARG ${_RULE_FLATCC_ARGS}) + if(_ARG STREQUAL "--reader") + list(APPEND _OUTS "${_SRC_FILENAME}_reader.h") + elseif(_ARG STREQUAL "--builder") + list(APPEND _OUTS "${_SRC_FILENAME}_builder.h") + elseif(_ARG STREQUAL "--verifier") + list(APPEND _OUTS "${_SRC_FILENAME}_verifier.h") + endif() + endforeach() + endforeach() + list(TRANSFORM _OUTS PREPEND "${CMAKE_CURRENT_BINARY_DIR}/") + + iree_get_executable_path(_FLATCC_BIN flatcc) + add_custom_command( + OUTPUT + ${_OUTS} + COMMAND + "${_FLATCC_BIN}" + -o "${CMAKE_CURRENT_BINARY_DIR}" + -I "${IREE_ROOT_DIR}" + ${_RULE_FLATCC_ARGS} + "${_RULE_SRCS}" + WORKING_DIRECTORY + "${CMAKE_CURRENT_SOURCE_DIR}" + MAIN_DEPENDENCY + ${_RULE_SRCS} + DEPENDS + ${_FLATCC_BIN} + ${_RULE_SRCS} + COMMAND_EXPAND_LISTS + ) + + set(_GEN_TARGET "${_NAME}_gen") + add_custom_target( + ${_GEN_TARGET} + DEPENDS + ${_OUTS} + ${_RULE_DEPS} + ) + + add_library(${_NAME} INTERFACE) + add_dependencies(${_NAME} ${_GEN_TARGET}) + target_include_directories(${_NAME} + INTERFACE + "$<BUILD_INTERFACE:${IREE_COMMON_INCLUDE_DIRS}>" + ${CMAKE_CURRENT_BINARY_DIR} + ) + target_link_libraries(${_NAME} + INTERFACE + flatcc::runtime + ${_RULE_LINKOPTS} + ${IREE_DEFAULT_LINKOPTS} + ) + target_compile_definitions(${_NAME} + INTERFACE + ${_RULE_DEFINES} + ) + target_compile_options(${_NAME} + INTERFACE + "-I${IREE_ROOT_DIR}/third_party/flatcc/include/flatcc/reflection/" + ) + + # Alias the iree_package_name library to iree::package::name. + # This lets us more clearly map to Bazel and makes it possible to + # disambiguate the underscores in paths vs. the separators. + iree_package_ns(_PACKAGE_NS) + add_library(${_PACKAGE_NS}::${_RULE_NAME} ALIAS ${_NAME}) +endfunction()
diff --git a/build_tools/cmake/flatbuffer_cc_library.cmake b/build_tools/cmake/flatbuffer_cc_library.cmake index febf234..86c38fc 100644 --- a/build_tools/cmake/flatbuffer_cc_library.cmake +++ b/build_tools/cmake/flatbuffer_cc_library.cmake
@@ -44,7 +44,7 @@ # NAME # base_schema # SRCS -# "a.cc" +# "a.fbs" # ) # flatbuffer_cc_library( # NAME @@ -52,11 +52,11 @@ # SRCS # "b.fbs" # DEPS -# iree::schemas::base_schema # not "awesome" ! +# iree::schemas::base_schema # PUBLIC # ) # -# flatbuffer_cc_library( +# iree_cc_binary( # NAME # main_lib # ...
diff --git a/build_tools/cmake/iree_check_test.cmake b/build_tools/cmake/iree_check_test.cmake index aa794bd..93fd934 100644 --- a/build_tools/cmake/iree_check_test.cmake +++ b/build_tools/cmake/iree_check_test.cmake
@@ -173,6 +173,17 @@ ${ARGN} ) + + string(TOUPPER ${_RULE_DRIVER} _UPPERCASE_DRIVER) + if(NOT IREE_HAL_DRIVER_${_UPPERCASE_DRIVER}) + return() + endif() + + string(TOUPPER ${_RULE_TARGET_BACKEND} _UPPERCASE_TARGET_BACKEND) + if(NOT IREE_TARGET_BACKEND_${_UPPERCASE_TARGET_BACKEND}) + return() + endif() + foreach(_SRC IN LISTS _RULE_SRCS) set(_TEST_NAME "${_RULE_NAME}_${_SRC}") iree_check_test(
diff --git a/build_tools/cmake/iree_copts.cmake b/build_tools/cmake/iree_copts.cmake index 7b18183..c39db21 100644 --- a/build_tools/cmake/iree_copts.cmake +++ b/build_tools/cmake/iree_copts.cmake
@@ -183,10 +183,19 @@ list(APPEND IREE_DEFAULT_COPTS ${FLATBUFFERS_COPTS}) #------------------------------------------------------------------------------- -# Third party: glslang +# Third party: flatcc #------------------------------------------------------------------------------- -set(ENABLE_CTEST OFF CACHE BOOL "" FORCE) +set(FLATCC_TEST OFF CACHE BOOL "" FORCE) +set(FLATCC_CXX_TEST OFF CACHE BOOL "" FORCE) +set(FLATCC_REFLECTION OFF CACHE BOOL "" FORCE) +set(FLATCC_ALLOW_WERROR OFF CACHE BOOL "" FORCE) + +if(CMAKE_CROSSCOMPILING) + set(FLATCC_RTONLY ON CACHE BOOL "" FORCE) +else() + set(FLATCC_RTONLY OFF CACHE BOOL "" FORCE) +endif() #------------------------------------------------------------------------------- # Third party: gtest
diff --git a/build_tools/cmake/iree_multipy.cmake b/build_tools/cmake/iree_multipy.cmake new file mode 100644 index 0000000..a7fd1f1 --- /dev/null +++ b/build_tools/cmake/iree_multipy.cmake
@@ -0,0 +1,462 @@ +# Copyright 2020 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +include(CMakeParseArguments) + +############################################################################### +# Configuration +############################################################################### + +function(iree_multipy_configure) + # Configure the defaults. + # Note that this is using the pybind11 configuration vars, which creates + # a fragile dependency. It would be better to derive these locally. + if(PYTHONLIBS_FOUND) + set(IREE_MULTIPY_DEFAULT_EXECUTABLE "${PYTHON_EXECUTABLE}" CACHE INTERNAL "Python executable" ) + set(IREE_MULTIPY_DEFAULT_INCLUDE_DIRS "${PYTHON_INCLUDE_DIRS}" CACHE INTERNAL "Python include dirs" ) + set(IREE_MULTIPY_DEFAULT_LIBRARIES "${PYTHON_LIBRARIES}" CACHE INTERNAL "Python libraries") + set(IREE_MULTIPY_DEFAULT_PREFIX "${PYTHON_MODULE_PREFIX}" CACHE INTERNAL "Python module prefix") + set(IREE_MULTIPY_DEFAULT_SUFFIX "${PYTHON_MODULE_SUFFIX}" CACHE INTERNAL "Python module suffix") + set(IREE_MULTIPY_DEFAULT_EXTENSION "${PYTHON_MODULE_EXTENSION}" CACHE INTERNAL "Python module extension") + endif() + + if(IREE_MULTIPY_VERSIONS) + set(IREE_MULTIPY_VERSIONS_EFFECTIVE "${IREE_MULTIPY_VERSIONS}" CACHE INTERNAL "Python extension versions") + else() + message(STATUS "Multi-python extension versions not found: using defaults") + set(IREE_MULTIPY_VERSIONS_EFFECTIVE "DEFAULT" CACHE INTERNAL "Python extension versions") + endif() + + # Report the multipy config. + message(STATUS "Multipy extension versions: ${IREE_MULTIPY_VERSIONS_EFFECTIVE}") + foreach(V ${IREE_MULTIPY_VERSIONS_EFFECTIVE}) + message(STATUS " - Multipy version ${V}") + message(STATUS " : EXECUTABLE = ${IREE_MULTIPY_${V}_EXECUTABLE}") + message(STATUS " : INCLUDE_DIRS = ${IREE_MULTIPY_${V}_INCLUDE_DIRS}") + message(STATUS " : LIBRARIES = ${IREE_MULTIPY_${V}_LIBRARIES}") + message(STATUS " : PREFIX = ${IREE_MULTIPY_${V}_PREFIX}") + message(STATUS " : SUFFIX = ${IREE_MULTIPY_${V}_SUFFIX}") + message(STATUS " : EXTENSION = ${IREE_MULTIPY_${V}_EXTENSION}") + + # Check for required settings. + if(NOT IREE_MULTIPY_${V}_INCLUDE_DIRS) + message(FATAL " MULTIPY version ${V}: No IREE_MULTIPY_${VER}_EXECUTABLE var") + endif() + if(NOT IREE_MULTIPY_${V}_INCLUDE_DIRS) + message(FATAL " MULTIPY version ${V}: No IREE_MULTIPY_${VER}_INCLUDE_DIRS var") + endif() + if(NOT IREE_MULTIPY_${V}_EXTENSION) + message(FATAL " MULTIPY version ${V}: No IREE_MULTIPY_${VER}_EXTENSION var") + endif() + endforeach() +endfunction() + +macro(_setup_iree_pyext_names) + iree_package_ns(_PACKAGE_NS) + # Replace dependencies passed by ::name with ::iree::package::name + list(TRANSFORM ARG_DEPS REPLACE "^::" "${_PACKAGE_NS}::") + list(TRANSFORM ARG_PYEXT_DEPS REPLACE "^::" "${_PACKAGE_NS}::") + # Prefix the library with the package name, so we get: iree_package_name. + iree_package_name(_PACKAGE_NAME) + set(_NAME "${_PACKAGE_NAME}_${ARG_NAME}") +endmacro() + +macro(_alias_iree_pyext_library declared_name version target) + # Alias the iree_package_name library to iree::package::name. + # This lets us more clearly map to Bazel and makes it possible to + # disambiguate the underscores in paths vs. the separators. + add_library(${_PACKAGE_NS}::${ARG_NAME}__${version} ALIAS ${target}) + iree_package_dir(_PACKAGE_DIR) +endmacro() + +############################################################################### +# Main user rules +############################################################################### + +# iree_pyext_module() +# +# Builds a native python module (.so/.dylib/.pyd). +# +# Parameters: +# NAME: name of target +# MODULE_NAME: Base-name of the module. +# SRCS: List of source files for the library +# COPTS: C options +# DEPS: List of other targets the test python libraries require +# PYEXT_DEPS: List of deps of extensions built with iree_pyext_(library|module) +function(iree_pyext_module) + cmake_parse_arguments(ARG + "" + "NAME;MODULE_NAME" + "SRCS;COPTS;DEPS;PYEXT_DEPS" + ${ARGN}) + _setup_iree_pyext_names() + + add_custom_target(${_NAME}) + + foreach(V ${IREE_MULTIPY_VERSIONS_EFFECTIVE}) + set(VER_NAME "${_NAME}__${V}") + + # If configured to link against libraries, build in SHARED mode (which + # disallows undefined symbols). Otherwise, build in MODULE mode, which + # does not enforce that. This should naturally do the right thing on + # each platform based on whether configured with a list of libraries to + # link or not. + set(LIBRARY_TYPE MODULE) + if(IREE_MULTIPY_${V}_LIBRARIES) + set(LIBRARY_TYPE SHARED) + endif() + + add_library(${VER_NAME} ${LIBRARY_TYPE} ${ARG_SRCS}) + add_dependencies(${_NAME} ${VER_NAME}) + set_target_properties( + ${VER_NAME} PROPERTIES + OUTPUT_NAME "${ARG_MODULE_NAME}" + PREFIX "${IREE_MULTIPY_${V}_PREFIX}" + SUFFIX "${IREE_MULTIPY_${V}_SUFFIX}${IREE_MULTIPY_${V}_EXTENSION}" + ) + + iree_pyext_pybind11_options(${VER_NAME}) + target_include_directories(${VER_NAME} + PUBLIC + "${IREE_MULTIPY_${V}_INCLUDE_DIRS}" + "$<BUILD_INTERFACE:${IREE_COMMON_INCLUDE_DIRS}>" + ) + target_link_libraries(${VER_NAME} + PRIVATE + ${IREE_DEFAULT_LINKOPTS} + ${IREE_MULTIPY_${V}_LIBRARIES} + ) + target_compile_options(${VER_NAME} + INTERFACE + ${IREE_DEFAULT_COPTS} + PRIVATE + ${ARG_COPTS} + ) + + # Defer computing transitive dependencies and calling target_link_libraries() + # until all libraries have been declared. + # Track target and deps, use in iree_complete_py_extension_link_options() later. + # See iree_complete_py_extension_link_options() in iree_py_extension.cmake + # TODO: Move that implementation here. + set(TRANSFORMED_PYEXT_DEPS "${ARG_PYEXT_DEPS}") + list(TRANSFORM TRANSFORMED_PYEXT_DEPS APPEND "__${V}") + set_property(GLOBAL APPEND PROPERTY _IREE_PY_EXTENSION_NAMES "${VER_NAME}") + set_property(TARGET ${VER_NAME} PROPERTY DIRECT_DEPS ${ARG_DEPS} ${TRANSFORMED_PYEXT_DEPS}) + _alias_iree_pyext_library("${ARG_NAME}" "${V}" ${VER_NAME}) + endforeach() +endfunction() + +# iree_pyext_library() +# +# Builds a C++ library to be included in an iree_pyext_module. +# +# Parameters: +# NAME: name of target +# SRCS: List of source files for the library +# COPTS: C options +# DEPS: List of other targets the test python libraries require +# PYEXT_DEPS: List of deps of extensions built with iree_pyext_(library|module) +function(iree_pyext_library) + cmake_parse_arguments(ARG + "" + "NAME" + "SRCS;COPTS;DEPS;PYEXT_DEPS" + ${ARGN}) + _setup_iree_pyext_names() + + foreach(V ${IREE_MULTIPY_VERSIONS_EFFECTIVE}) + set(VER_NAME "${_NAME}__${V}") + add_library(${VER_NAME} STATIC ${ARG_SRCS}) + iree_pyext_pybind11_options(${VER_NAME}) + target_include_directories(${VER_NAME} + PUBLIC + "${IREE_MULTIPY_${V}_INCLUDE_DIRS}" + "$<BUILD_INTERFACE:${IREE_COMMON_INCLUDE_DIRS}>" + ) + set(TRANSFORMED_PYEXT_DEPS "${ARG_PYEXT_DEPS}") + list(TRANSFORM TRANSFORMED_PYEXT_DEPS APPEND "__${V}") + target_link_libraries(${VER_NAME} + PUBLIC + ${ARG_DEPS} + ${TRANSFORMED_PYEXT_DEPS} + PRIVATE + ${IREE_DEFAULT_LINKOPTS} + ) + target_compile_options(${VER_NAME} + INTERFACE + ${IREE_DEFAULT_COPTS} + PRIVATE + ${ARG_COPTS} + ) + _alias_iree_pyext_library("${ARG_NAME}" "${V}" ${VER_NAME}) + endforeach() +endfunction() + +# iree_py_library() +# +# CMake function to imitate Bazel's iree_py_library rule. +# +# Parameters: +# NAME: name of target +# SRCS: List of source files for the library +# DEPS: List of other targets the test python libraries require +# PYEXT_DEPS: List of deps of extensions built with iree_pyext_module +function(iree_py_library) + cmake_parse_arguments( + ARG + "" + "NAME" + "SRCS;DEPS;PYEXT_DEPS" + ${ARGN} + ) + + iree_package_ns(_PACKAGE_NS) + # Replace dependencies passed by ::name with ::iree::package::name + list(TRANSFORM ARG_DEPS REPLACE "^::" "${_PACKAGE_NS}::") + + iree_package_name(_PACKAGE_NAME) + set(_NAME "${_PACKAGE_NAME}_${ARG_NAME}") + + # Add path to each source file + list(TRANSFORM ARG_SRCS PREPEND "${CMAKE_CURRENT_SOURCE_DIR}/") + + add_custom_target(${_NAME} ALL + COMMAND ${CMAKE_COMMAND} -E copy ${ARG_SRCS} "${CMAKE_CURRENT_BINARY_DIR}/" + DEPENDS ${ARG_DEPS} + ) + + # Add PYEXT_DEPS. + if(${ARG_PYEXT_DEPS}) + foreach(V ${IREE_MULTIPY_VERSIONS_EFFECTIVE}) + list(TRANSFORM ARG_PYEXT_DEPS APPEND "__${V}") + add_dependencies(${_NAME} ${ARG_PYEXT_DEPS}) + endforeach() + endif() +endfunction() + +function(iree_pyext_pybind11_options name) + target_include_directories(${name} + PRIVATE + ${PYBIND11_INCLUDE_DIR} + ) + target_compile_options(${name} + PRIVATE + $<$<OR:$<CXX_COMPILER_ID:Clang>,$<CXX_COMPILER_ID:AppleClang>,$<CXX_COMPILER_ID:GNU>>: + -frtti -fexceptions + # Noisy pybind warnings + -Wno-unused-value + -Wno-covered-switch-default + > + $<$<CXX_COMPILER_ID:MSVC>: + # Enable RTTI and exceptions. + /EHsc /GR> + ) + set_target_properties( + ${name} PROPERTIES CXX_VISIBILITY_PRESET "hidden") +endfunction() + +# iree_py_test() +# +# CMake function to imitate Bazel's iree_py_test rule. +# +# Parameters: +# NAME: name of test +# SRCS: List of source file +# DEPS: List of deps the test requires +# LABELS: Additional labels to apply to the test. The package path is added +# automatically. + +function(iree_py_test) + if(NOT IREE_BUILD_TESTS) + return() + endif() + + cmake_parse_arguments( + _RULE + "" + "NAME" + "SRCS;DEPS;LABELS" + ${ARGN} + ) + + iree_package_name(_PACKAGE_NAME) + set(_NAME "${_PACKAGE_NAME}_${_RULE_NAME}") + + iree_package_ns(_PACKAGE_NS) + string(REPLACE "::" "/" _PACKAGE_PATH ${_PACKAGE_NS}) + set(_NAME_PATH "${_PACKAGE_PATH}:${_RULE_NAME}") + list(APPEND _RULE_LABELS "${_PACKAGE_PATH}") + + foreach(V ${IREE_MULTIPY_VERSIONS_EFFECTIVE}) + set(VER_NAME "${_NAME_PATH}__${V}") + add_test( + NAME ${VER_NAME} + COMMAND + "${CMAKE_SOURCE_DIR}/build_tools/cmake/run_test.${IREE_HOST_SCRIPT_EXT}" + "${IREE_MULTIPY_${V}_EXECUTABLE}" + "${CMAKE_CURRENT_SOURCE_DIR}/${_RULE_SRCS}" + WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR} + ) + + set_property(TEST ${VER_NAME} PROPERTY LABELS "${_RULE_LABELS}") + set_property(TEST ${VER_NAME} PROPERTY ENVIRONMENT "PYTHONPATH=${CMAKE_BINARY_DIR}/bindings/python:$ENV{PYTHONPATH};TEST_TMPDIR=${_NAME}_${V}_test_tmpdir") + # TODO(marbre): Find out how to add deps to tests. + # Similar to _RULE_DATA in iree_lit_test(). + endforeach() +endfunction() + +############################################################################### +# Always-link/transitive dependency management +############################################################################### + +# Lists all transitive dependencies of DIRECT_DEPS in TRANSITIVE_DEPS. +function(_iree_transitive_dependencies DIRECT_DEPS TRANSITIVE_DEPS) + set(_TRANSITIVE "") + + foreach(_DEP ${DIRECT_DEPS}) + _iree_transitive_dependencies_helper(${_DEP} _TRANSITIVE) + endforeach(_DEP) + + set(${TRANSITIVE_DEPS} "${_TRANSITIVE}" PARENT_SCOPE) +endfunction() + +# Recursive helper function for _iree_transitive_dependencies. +# Performs a depth-first search through the dependency graph, appending all +# dependencies of TARGET to the TRANSITIVE_DEPS list. +function(_iree_transitive_dependencies_helper TARGET TRANSITIVE_DEPS) + if (NOT TARGET "${TARGET}") + # Excluded from the project, or invalid name? Just ignore. + return() + endif() + + # Resolve aliases, canonicalize name formatting. + get_target_property(_ALIASED_TARGET ${TARGET} ALIASED_TARGET) + if(_ALIASED_TARGET) + set(_TARGET_NAME ${_ALIASED_TARGET}) + else() + string(REPLACE "::" "_" _TARGET_NAME ${TARGET}) + endif() + + set(_RESULT "${${TRANSITIVE_DEPS}}") + if (${_TARGET_NAME} IN_LIST _RESULT) + # Already visited, ignore. + return() + endif() + + # Append this target to the list. Dependencies of this target will be added + # (if valid and not already visited) in recursive function calls. + list(APPEND _RESULT ${_TARGET_NAME}) + + # Check for non-target identifiers again after resolving the alias. + if (NOT TARGET ${_TARGET_NAME}) + return() + endif() + + # Get the list of direct dependencies for this target. + get_target_property(_TARGET_TYPE ${_TARGET_NAME} TYPE) + if(NOT ${_TARGET_TYPE} STREQUAL "INTERFACE_LIBRARY") + get_target_property(_TARGET_DEPS ${_TARGET_NAME} LINK_LIBRARIES) + else() + get_target_property(_TARGET_DEPS ${_TARGET_NAME} INTERFACE_LINK_LIBRARIES) + endif() + + if(_TARGET_DEPS) + # Recurse on each dependency. + foreach(_TARGET_DEP ${_TARGET_DEPS}) + _iree_transitive_dependencies_helper(${_TARGET_DEP} _RESULT) + endforeach(_TARGET_DEP) + endif() + + # Propagate the augmented list up to the parent scope. + set(${TRANSITIVE_DEPS} "${_RESULT}" PARENT_SCOPE) +endfunction() + +# Sets target_link_libraries() on all registered py extensions. +# This must be called after all libraries have been declared. +function(iree_complete_py_extension_link_options) + get_property(_NAMES GLOBAL PROPERTY _IREE_PY_EXTENSION_NAMES) + + foreach(_NAME ${_NAMES}) + get_target_property(_DIRECT_DEPS ${_NAME} DIRECT_DEPS) + + # List all dependencies, including transitive dependencies, then split the + # dependency list into one for whole archive (ALWAYSLINK) and one for + # standard linking (which only links in symbols that are directly used). + _iree_transitive_dependencies("${_DIRECT_DEPS}" _TRANSITIVE_DEPS) + set(_ALWAYS_LINK_DEPS "") + set(_STANDARD_DEPS "") + foreach(_DEP ${_TRANSITIVE_DEPS}) + # Check if _DEP is a library with the ALWAYSLINK property set. + set(_DEP_IS_ALWAYSLINK OFF) + if (TARGET ${_DEP}) + get_target_property(_DEP_TYPE ${_DEP} TYPE) + if(${_DEP_TYPE} STREQUAL "INTERFACE_LIBRARY") + # Can't be ALWAYSLINK since it's an INTERFACE library. + # We also can't even query for the property, since it isn't allowlisted. + else() + get_target_property(_DEP_IS_ALWAYSLINK ${_DEP} ALWAYSLINK) + endif() + endif() + + # Append to the corresponding list of deps. + if(_DEP_IS_ALWAYSLINK) + list(APPEND _ALWAYS_LINK_DEPS ${_DEP}) + + # For MSVC, also add a `-WHOLEARCHIVE:` version of the dep. + # CMake treats -WHOLEARCHIVE[:lib] as a link flag and will not actually + # try to link the library in, so we need the flag *and* the dependency. + # For macOS, also add a `-Wl,-force_load` version of the dep. + if(MSVC) + get_target_property(_ALIASED_TARGET ${_DEP} ALIASED_TARGET) + if (_ALIASED_TARGET) + list(APPEND _ALWAYS_LINK_DEPS "-WHOLEARCHIVE:${_ALIASED_TARGET}") + else() + list(APPEND _ALWAYS_LINK_DEPS "-WHOLEARCHIVE:${_DEP}") + endif() + elseif(APPLE) + get_target_property(_ALIASED_TARGET ${_DEP} ALIASED_TARGET) + if (_ALIASED_TARGET) + list(APPEND _ALWAYS_LINK_DEPS "-Wl,-force_load $<TARGET_FILE:${_ALIASED_TARGET}>") + else() + list(APPEND _ALWAYS_LINK_DEPS "-Wl,-force_load $<TARGET_FILE:${_DEP}>") + endif() + endif() + else() + list(APPEND _STANDARD_DEPS ${_DEP}) + endif() + endforeach(_DEP) + + # Call into target_link_libraries with the lists of deps. + if(MSVC OR APPLE) + target_link_libraries(${_NAME} + PUBLIC + ${_ALWAYS_LINK_DEPS} + ${_STANDARD_DEPS} + PRIVATE + ${_RULE_LINKOPTS} + ) + else() + target_link_libraries(${_NAME} + PUBLIC + "-Wl,--whole-archive" + ${_ALWAYS_LINK_DEPS} + "-Wl,--no-whole-archive" + ${_STANDARD_DEPS} + PRIVATE + ${_RULE_LINKOPTS} + ) + endif() + endforeach(_NAME) +endfunction()
diff --git a/build_tools/cmake/run_android_test.sh b/build_tools/cmake/run_android_test.sh index 292890c..0c7510d 100755 --- a/build_tools/cmake/run_android_test.sh +++ b/build_tools/cmake/run_android_test.sh
@@ -35,7 +35,7 @@ set -x set -e -adb push $TEST_EXECUTABLE $TEST_ANDROID_ABS_DIR/$(basename $TEST_EXECUTABLE) +adb push $TEST_EXECUTABLE $TEST_ANDROID_ABS_DIR/$(basename $TEST_EXECUTABLE) 1>/dev/null if [ -n "$TEST_DATA" ]; then adb push $TEST_DATA $TEST_ANDROID_ABS_DIR/$(basename $TEST_DATA)
diff --git a/build_tools/docker/bazel/Dockerfile b/build_tools/docker/bazel/Dockerfile index e9a587c..c42521a 100644 --- a/build_tools/docker/bazel/Dockerfile +++ b/build_tools/docker/bazel/Dockerfile
@@ -27,40 +27,49 @@ FROM ubuntu:18.04 WORKDIR /usr/src/iree -RUN apt-get update - # Set environment variables. ENV CXX clang++ ENV CC clang ENV PYTHON_BIN /usr/bin/python3 ENV IREE_LLVMAOT_LINKER_PATH /usr/bin/ld -# Install git for updating IREE's submodules. -RUN apt-get install -y git +RUN apt-get update \ + && apt-get install -y \ + # git for updating IREE's submodules. + git \ + # utilities for later installations + unzip \ + zip \ + wget \ + # core IREE dependencies. + clang \ + libsdl2-dev + +# Disable apt-key parse waring. If someone knows how to do whatever the "proper" +# thing is then feel free. The warning complains about parsing apt-key output, +# which we're not even doing. +ARG APT_KEY_DONT_WARN_ON_DANGEROUS_USAGE=1 # Install Bazel. # https://docs.bazel.build/versions/master/install-ubuntu.html -ARG BAZEL_VERSION=2.1.0 -RUN apt-get install -y unzip zip wget \ - && wget "https://github.com/bazelbuild/bazel/releases/download/${BAZEL_VERSION?}/bazel-${BAZEL_VERSION?}-installer-linux-x86_64.sh" \ - && chmod +x "bazel-${BAZEL_VERSION?}-installer-linux-x86_64.sh" \ - && "./bazel-${BAZEL_VERSION?}-installer-linux-x86_64.sh" --user \ - && rm "bazel-${BAZEL_VERSION?}-installer-linux-x86_64.sh" -# Install a newer version of Bazel. We don't need the full installation now. -# Just need to provide a different version for the version-identification -# wrapper script to find in /root/.bazel/bin +ARG BAZEL_VERSION=3.3.1 +# Change to a new version if upgrading Bazel. ARG NEW_BAZEL_VERSION=3.3.1 -RUN cd "/root/.bazel/bin" \ - && wget "https://releases.bazel.build/${NEW_BAZEL_VERSION?}/release/bazel-${NEW_BAZEL_VERSION?}-linux-x86_64" \ - && chmod +x "bazel-${NEW_BAZEL_VERSION?}-linux-x86_64" -# ENV does not allow ${variable?} syntax. -ENV PATH "/root/bin:${PATH}" +RUN wget -qO - https://bazel.build/bazel-release.pub.gpg | apt-key add - \ + && echo "deb [arch=amd64] https://storage.googleapis.com/bazel-apt stable jdk1.8" \ + | tee /etc/apt/sources.list.d/bazel.list \ + && apt-get update \ + # Install Bazel pinned at the version we want. Optionally install an + # additional version of Bazel to ease upgrades (modify NEW_BAZEL_VERSION + # above). Bazel does some shenanigans to select the correct version based on + # your .bazelversion file. When upgrading, we therefore need to have both the + # old and new version. When the versions are the same this second installation + # is effectively a noop. + && apt-get install "bazel=${BAZEL_VERSION?}" "bazel-${NEW_BAZEL_VERSION?}" -# Install core IREE dependencies. -RUN apt-get install -y clang libsdl2-dev - -# Install python2 numpy. Temporary fix for issue #1737: -# https://github.com/google/iree/issues/1737 -RUN apt-get install -y python-pip \ - && python -m pip install --upgrade pip \ - && python -m pip install numpy +# TF requires python2 numpy at configure time... +# TODO(#1737): Remove this +RUN apt-get update \ + && apt-get install -y python-pip \ + && python -m pip install --upgrade pip \ + && python -m pip install numpy
diff --git a/build_tools/docker/bazel_bindings/Dockerfile b/build_tools/docker/bazel_bindings/Dockerfile index f7ef30f..8f07958 100644 --- a/build_tools/docker/bazel_bindings/Dockerfile +++ b/build_tools/docker/bazel_bindings/Dockerfile
@@ -27,6 +27,11 @@ FROM gcr.io/iree-oss/bazel # Install python3 and numpy. -RUN apt-get install -y python3 python3-dev python3-pip python3-setuptools \ +RUN apt-get update \ + && apt-get install -y \ + python3 \ + python3-dev \ + python3-pip \ + python3-setuptools \ && python3 -m pip install --upgrade pip \ && python3 -m pip install numpy
diff --git a/build_tools/docker/bazel_tensorflow/Dockerfile b/build_tools/docker/bazel_tensorflow/Dockerfile index a117990..0c37354 100644 --- a/build_tools/docker/bazel_tensorflow/Dockerfile +++ b/build_tools/docker/bazel_tensorflow/Dockerfile
@@ -24,9 +24,7 @@ # gcr.io/iree-oss/bazel-tensorflow # Set up the image and working directory. -FROM gcr.io/iree-oss/bazel +FROM gcr.io/iree-oss/bazel-bindings -# Install python3, tensorflow and numpy. -RUN apt-get install -y python3 python3-dev python3-pip python3-setuptools \ - && python3 -m pip install --upgrade pip \ - && python3 -m pip install numpy tf-nightly +# Install tensorflow. +RUN python3 -m pip install tf-nightly
diff --git a/build_tools/docker/build_and_update_gcr.py b/build_tools/docker/build_and_update_gcr.py index 90c7454..531088f 100755 --- a/build_tools/docker/build_and_update_gcr.py +++ b/build_tools/docker/build_and_update_gcr.py
@@ -20,30 +20,33 @@ """ import argparse +import functools import os import subprocess +import sys IREE_GCR_URL = 'gcr.io/iree-oss/' DOCKER_DIR = 'build_tools/docker/' -IMAGES = [ - 'bazel', - 'bazel-bindings', - 'bazel-tensorflow', - 'cmake', - 'cmake-android', - 'cmake-nvidia', - 'rbe-toolchain', -] -IMAGES_HELP = [f'`{name}`' for name in IMAGES] -IMAGES_HELP = f'{", ".join(IMAGES_HELP[:-1])} or {IMAGES_HELP[-1]}' - -# Map from image names to images that depend on them. -IMAGES_TO_DEPENDENT_IMAGES = { - 'bazel': ['bazel-bindings', 'bazel-tensorflow'], - 'cmake': ['cmake-android', 'cmake-nvidia'], +# Map from image names to images that they depend on. +IMAGES_TO_DEPENDENCIES = { + 'bazel': [], + 'bazel-bindings': ['bazel'], + 'bazel-tensorflow': ['bazel-bindings'], + 'cmake': [], + 'cmake-android': ['cmake'], + 'cmake-nvidia': ['cmake'], + 'rbe-toolchain': [], } +IMAGES_TO_DEPENDENT_IMAGES = {k: [] for k in IMAGES_TO_DEPENDENCIES.keys()} +for image, dependencies in IMAGES_TO_DEPENDENCIES.items(): + for dependency in dependencies: + IMAGES_TO_DEPENDENT_IMAGES[dependency].append(image) + +IMAGES_HELP = [f'`{name}`' for name in IMAGES_TO_DEPENDENCIES.keys()] +IMAGES_HELP = f'{", ".join(IMAGES_HELP)} or `all`' + RBE_MESSAGE = """ Remember to update the `rbe_default` digest in the `WORKSPACE` file to reflect the new digest for the container. @@ -57,8 +60,10 @@ description="Build IREE's Docker images and optionally push them to GCR.") parser.add_argument( '--image', + dest='images', type=str, required=True, + action='append', help=f'Name of the image to build: {IMAGES_HELP}.') parser.add_argument( '--tag', @@ -73,34 +78,77 @@ help='Push the built images to GCR. Requires gcloud authorization.') args = parser.parse_args() - if args.image not in IMAGES: - raise parser.error('Expected --image to be one of:\n' - f' {IMAGES_HELP}\n' - f'but got `{args.image}`.') - + for image in args.images: + if image == 'all': + args.images = IMAGES_TO_DEPENDENCIES.keys() + elif image not in IMAGES_TO_DEPENDENCIES.keys(): + raise parser.error('Expected --image to be one of:\n' + f' {IMAGES_HELP}\n' + f'but got `{image}`.') return args +def cmp_images_by_dependency(image1, image2): + if image2 in IMAGES_TO_DEPENDENT_IMAGES[image1]: + return -1 + if image1 in IMAGES_TO_DEPENDENT_IMAGES[image2]: + return 1 + return (image1 > image2) - (image1 < image2) + + +def run_command(command): + print(f'Running: {" ".join(command)}') + process = subprocess.Popen( + command, + bufsize=1, + stderr=subprocess.STDOUT, + stdout=subprocess.PIPE, + text=True) + for line in process.stdout: + print(line, end='') + + return process.poll() + + +def check_command(command): + exit_code = run_command(command) + if exit_code != 0: + print(f'Command failed: {" ".join(command)}') + sys.exit(exit_code) + + if __name__ == '__main__': args = parse_arguments() # Ensure the user has the correct authorization if they try to push to GCR. if args.push: - subprocess.check_output(['gcloud', 'auth', 'configure-docker']) + if run_command(['which', 'gcloud']) != 0: + print('gcloud not found.' + ' See https://cloud.google.com/sdk/install for installation.') + sys.exit(1) + check_command(['gcloud', 'auth', 'configure-docker']) - # Check if any images depend on `args.image` and update them if they do. - images_to_update = [args.image] - if args.image in IMAGES_TO_DEPENDENT_IMAGES: - images_to_update.extend(IMAGES_TO_DEPENDENT_IMAGES[args.image]) + # Check if any images depend on `args.images` and update them if they do. + images_to_update_set = set() + to_check = list(args.images) + while to_check: + image = to_check.pop() + if image not in images_to_update_set: + images_to_update_set.add(image) + to_check.extend(IMAGES_TO_DEPENDENT_IMAGES[image]) + # Topo sort by image dependency + images_to_update = sorted( + images_to_update_set, key=functools.cmp_to_key(cmp_images_by_dependency)) + + print(f'Also updating dependent images. Will update: {images_to_update}') for image in images_to_update: print(f'Updating image {image}') image_url = os.path.join(IREE_GCR_URL, f'{image}:{args.tag}') image_path = os.path.join(DOCKER_DIR, image.replace('-', '_')) - subprocess.check_output( - ['docker', 'build', '--tag', image_url, image_path]) + check_command(['docker', 'build', '--tag', image_url, image_path]) if args.push: - subprocess.check_output(['docker', 'push', image_url]) + check_command(['docker', 'push', image_url]) if 'rbe-toolchain' in images_to_update: print(RBE_MESSAGE)
diff --git a/build_tools/docker/cmake/Dockerfile b/build_tools/docker/cmake/Dockerfile index 92cece8..bde1f4d 100644 --- a/build_tools/docker/cmake/Dockerfile +++ b/build_tools/docker/cmake/Dockerfile
@@ -27,14 +27,25 @@ FROM ubuntu:18.04 WORKDIR /usr/src/iree/ -RUN apt-get update -# TODO: Remove this if the `apt-get install` below works without it again. -RUN apt update +RUN apt-get update \ + && apt-get install -y \ + # git for updating IREE's submodules. + git \ + # For later installations + wget \ + # For building with ninja + ninja-build \ + # For bootstrapping the cmake installation + cmake \ + # core IREE dependencies. + clang \ + libsdl2-dev \ + libssl-dev # Update cmake to v3.13+, which is ahead of apt-get's version (3.10.2). # Install dependencies, including an old version of cmake to bootstrap. ENV CMAKE_VERSION 3.13.5 -RUN apt-get install -y clang cmake libssl-dev wget \ +RUN apt-get update \ && mkdir ./cmake_install \ && cd cmake_install \ && wget "https://github.com/Kitware/CMake/releases/download/v${CMAKE_VERSION?}/cmake-${CMAKE_VERSION?}.tar.gz" \ @@ -44,14 +55,12 @@ && make \ && make install -# Install dependencies. -RUN apt-get install -y \ - git \ - ninja-build \ - python3 \ - python3-pip \ - python3-setuptools \ - # Install dependencies for the python bindings tests. +# Dependencies for the python bindings tests. +RUN apt-get update \ + && apt-get install -y \ + python3 \ + python3-pip \ + python3-setuptools \ && python3 -m pip install --upgrade pip \ && python3 -m pip install numpy absl-py
diff --git a/build_tools/docker/cmake_nvidia/Dockerfile b/build_tools/docker/cmake_nvidia/Dockerfile index 8679656..c25519f 100644 --- a/build_tools/docker/cmake_nvidia/Dockerfile +++ b/build_tools/docker/cmake_nvidia/Dockerfile
@@ -40,13 +40,15 @@ ARG VULKAN_SDK_VERSION=1.2.141 -# Disable apt-key parse waring. -ENV APT_KEY_DONT_WARN_ON_DANGEROUS_USAGE=1 +# Disable apt-key parse waring. If someone knows how to do whatever the "proper" +# thing is then feel free. The warning complains about parsing apt-key output, +# which we're not even doing. +ARG APT_KEY_DONT_WARN_ON_DANGEROUS_USAGE=1 -# Disable prompt during keyboard configuration. -ENV DEBIAN_FRONTEND=noninteractive - -RUN wget -qO - http://packages.lunarg.com/lunarg-signing-key-pub.asc | apt-key add - \ - && wget -qO /etc/apt/sources.list.d/lunarg-vulkan-$VULKAN_SDK_VERSION-bionic.list http://packages.lunarg.com/vulkan/$VULKAN_SDK_VERSION/lunarg-vulkan-$VULKAN_SDK_VERSION-bionic.list \ +RUN wget -qO - http://packages.lunarg.com/lunarg-signing-key-pub.asc \ + | apt-key add - \ + && wget -qO \ + "/etc/apt/sources.list.d/lunarg-vulkan-${VULKAN_SDK_VERSION?}-bionic.list" \ + "http://packages.lunarg.com/vulkan/${VULKAN_SDK_VERSION?}/lunarg-vulkan-${VULKAN_SDK_VERSION?}-bionic.list" \ && apt-get update \ && apt-get install -y vulkan-sdk nvidia-driver-440
diff --git a/build_tools/docker/rbe_toolchain/Dockerfile b/build_tools/docker/rbe_toolchain/Dockerfile index c2df299..878aca3 100755 --- a/build_tools/docker/rbe_toolchain/Dockerfile +++ b/build_tools/docker/rbe_toolchain/Dockerfile
@@ -21,23 +21,29 @@ FROM gcr.io/cloud-marketplace/google/rbe-ubuntu16-04@sha256:1a8ed713f40267bb51fe17de012fa631a20c52df818ccb317aaed2ee068dfc61 -RUN apt-get update -RUN apt-get install -y python3 python3-pip -RUN python3 -m pip install --upgrade pip -RUN python3 -m pip install numpy +RUN apt-get update \ + && apt-get install -y \ + python3 \ + python3-pip \ + && python3 -m pip install --upgrade pip \ + && python3 -m pip install numpy -# Install dependencies for python3.6-dev -RUN apt-get install -y software-properties-common +# Dependency for python3.6-dev. Needs to be installed separately from the above +# for... some reason +RUN apt-get update && apt-get install -y software-properties-common + # apt-add-repository requires a version of python with the softwareproperties # module. To use this command, we: # 1. remove the symlink to python3 from python3.6 and symlink it to python3.5 # 2. run apt-add-repository with python3 = python3.5 # 3. resymlink python3 to /opt/python3.6/bin/python3.6 # See https://github.com/google/iree/issues/1966 for more information. -RUN rm /usr/bin/python3 && ln -s /usr/bin/python3.5 /usr/bin/python3 -RUN add-apt-repository ppa:deadsnakes/ppa -RUN rm /usr/bin/python3 && ln -s /opt/python3.6/bin/python3.6 /usr/bin/python3 +RUN rm /usr/bin/python3 \ + && ln -s /usr/bin/python3.5 /usr/bin/python3 \ + && add-apt-repository ppa:deadsnakes/ppa \ + && rm /usr/bin/python3 \ + && ln -s /opt/python3.6/bin/python3.6 /usr/bin/python3 # Install python3.6-dev -RUN apt-get update -RUN apt-get install -y python3.6 python3.6-dev +RUN apt-get update \ + && apt-get install -y python3.6 python3.6-dev
diff --git a/kokoro/gcp_ubuntu/bazel/bindings/build.sh b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/bindings/build.sh similarity index 100% rename from kokoro/gcp_ubuntu/bazel/bindings/build.sh rename to build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/bindings/build.sh
diff --git a/kokoro/gcp_ubuntu/bazel/bindings/build_kokoro.sh b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/bindings/build_kokoro.sh old mode 100644 new mode 100755 similarity index 73% rename from kokoro/gcp_ubuntu/bazel/bindings/build_kokoro.sh rename to build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/bindings/build_kokoro.sh index e2a9bd3..24cbabf --- a/kokoro/gcp_ubuntu/bazel/bindings/build_kokoro.sh +++ b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/bindings/build_kokoro.sh
@@ -16,28 +16,28 @@ # Build and test IREE's bindings within the gcr.io/iree-oss/bazel-bindings # image using Kokoro. +# Requires the environment variables KOKORO_ROOT and KOKORO_ARTIFACTS_DIR, which +# are set by Kokoro. -set -e set -x +set -e +set -o pipefail # Print the UTC time when set -x is on export PS4='[$(date -u "+%T %Z")] ' -# Kokoro checks out the repository here. -WORKDIR="${KOKORO_ARTIFACTS_DIR?}/github/iree" +source "${KOKORO_ARTIFACTS_DIR?}/github/iree/build_tools/kokoro/gcp_ubuntu/docker_common.sh" -# Mount the checked out repository, make that the working directory and run the -# tests in the bazel-bindings image. -docker run \ - --volume "${WORKDIR?}:${WORKDIR?}" \ - --workdir="${WORKDIR?}" \ - --rm \ +# Sets DOCKER_RUN_ARGS +docker_setup + +docker run "${DOCKER_RUN_ARGS[@]?}" \ gcr.io/iree-oss/bazel-bindings:prod \ - kokoro/gcp_ubuntu/bazel/bindings/build.sh + build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/bindings/build.sh # Kokoro will rsync this entire directory back to the executor orchestrating the # build which takes forever and is totally useless. -sudo rm -rf "${KOKORO_ARTIFACTS_DIR?}"/* +rm -rf "${KOKORO_ARTIFACTS_DIR?}"/* # Print out artifacts dir contents after deleting them as a coherence check. ls -1a "${KOKORO_ARTIFACTS_DIR?}/"
diff --git a/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/bindings/common.cfg b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/bindings/common.cfg new file mode 100644 index 0000000..8a49430 --- /dev/null +++ b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/bindings/common.cfg
@@ -0,0 +1,20 @@ +# Format: //devtools/kokoro/config/proto/build.proto + +# Copyright 2019 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Common configuration for Kokoro builds that run the bindings build with bazel +# on linux. + +build_file: "iree/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/bindings/build_kokoro.sh"
diff --git a/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/bindings/continuous.cfg b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/bindings/continuous.cfg new file mode 100644 index 0000000..50a7eed --- /dev/null +++ b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/bindings/continuous.cfg
@@ -0,0 +1,19 @@ +# Format: //devtools/kokoro/config/proto/build.proto + +# Copyright 2019 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Deliberately blank as everything necessary is configured in common files, but +# file must still exist to match corresponding (Google internal) job +# configurations that trigger the builds.
diff --git a/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/bindings/google.cfg b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/bindings/google.cfg new file mode 100644 index 0000000..50a7eed --- /dev/null +++ b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/bindings/google.cfg
@@ -0,0 +1,19 @@ +# Format: //devtools/kokoro/config/proto/build.proto + +# Copyright 2019 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Deliberately blank as everything necessary is configured in common files, but +# file must still exist to match corresponding (Google internal) job +# configurations that trigger the builds.
diff --git a/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/bindings/main.cfg b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/bindings/main.cfg new file mode 100644 index 0000000..50a7eed --- /dev/null +++ b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/bindings/main.cfg
@@ -0,0 +1,19 @@ +# Format: //devtools/kokoro/config/proto/build.proto + +# Copyright 2019 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Deliberately blank as everything necessary is configured in common files, but +# file must still exist to match corresponding (Google internal) job +# configurations that trigger the builds.
diff --git a/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/bindings/presubmit.cfg b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/bindings/presubmit.cfg new file mode 100644 index 0000000..50a7eed --- /dev/null +++ b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/bindings/presubmit.cfg
@@ -0,0 +1,19 @@ +# Format: //devtools/kokoro/config/proto/build.proto + +# Copyright 2019 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Deliberately blank as everything necessary is configured in common files, but +# file must still exist to match corresponding (Google internal) job +# configurations that trigger the builds.
diff --git a/kokoro/gcp_ubuntu/bazel/core/build.sh b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/core/build.sh similarity index 100% rename from kokoro/gcp_ubuntu/bazel/core/build.sh rename to build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/core/build.sh
diff --git a/kokoro/gcp_ubuntu/bazel/core/build_kokoro.sh b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/core/build_kokoro.sh old mode 100644 new mode 100755 similarity index 72% rename from kokoro/gcp_ubuntu/bazel/core/build_kokoro.sh rename to build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/core/build_kokoro.sh index 031eab8..da07cdd --- a/kokoro/gcp_ubuntu/bazel/core/build_kokoro.sh +++ b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/core/build_kokoro.sh
@@ -16,28 +16,28 @@ # Build and test IREE's core within the gcr.io/iree-oss/bazel image using # Kokoro. +# Requires the environment variables KOKORO_ROOT and KOKORO_ARTIFACTS_DIR, which +# are set by Kokoro. -set -e set -x +set -e +set -o pipefail # Print the UTC time when set -x is on export PS4='[$(date -u "+%T %Z")] ' -# Kokoro checks out the repository here. -WORKDIR="${KOKORO_ARTIFACTS_DIR?}/github/iree" +source "${KOKORO_ARTIFACTS_DIR?}/github/iree/build_tools/kokoro/gcp_ubuntu/docker_common.sh" -# Mount the checked out repository, make that the working directory and run the -# tests in the bazel image. -docker run \ - --volume "${WORKDIR?}:${WORKDIR?}" \ - --workdir="${WORKDIR?}" \ - --rm \ +# Sets DOCKER_RUN_ARGS +docker_setup + +docker run "${DOCKER_RUN_ARGS[@]?}" \ gcr.io/iree-oss/bazel:prod \ - kokoro/gcp_ubuntu/bazel/core/build.sh + build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/core/build.sh # Kokoro will rsync this entire directory back to the executor orchestrating the # build which takes forever and is totally useless. -sudo rm -rf "${KOKORO_ARTIFACTS_DIR?}"/* +rm -rf "${KOKORO_ARTIFACTS_DIR?}"/* # Print out artifacts dir contents after deleting them as a coherence check. ls -1a "${KOKORO_ARTIFACTS_DIR?}/"
diff --git a/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/core/common.cfg b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/core/common.cfg new file mode 100755 index 0000000..3a22d10 --- /dev/null +++ b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/core/common.cfg
@@ -0,0 +1,20 @@ +# Format: //devtools/kokoro/config/proto/build.proto + +# Copyright 2019 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Common configuration for Kokoro builds that run the core build with bazel on +# linux. + +build_file: "iree/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/core/build_kokoro.sh"
diff --git a/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/core/continuous.cfg b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/core/continuous.cfg new file mode 100755 index 0000000..50a7eed --- /dev/null +++ b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/core/continuous.cfg
@@ -0,0 +1,19 @@ +# Format: //devtools/kokoro/config/proto/build.proto + +# Copyright 2019 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Deliberately blank as everything necessary is configured in common files, but +# file must still exist to match corresponding (Google internal) job +# configurations that trigger the builds.
diff --git a/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/core/google.cfg b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/core/google.cfg new file mode 100755 index 0000000..50a7eed --- /dev/null +++ b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/core/google.cfg
@@ -0,0 +1,19 @@ +# Format: //devtools/kokoro/config/proto/build.proto + +# Copyright 2019 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Deliberately blank as everything necessary is configured in common files, but +# file must still exist to match corresponding (Google internal) job +# configurations that trigger the builds.
diff --git a/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/core/main.cfg b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/core/main.cfg new file mode 100755 index 0000000..50a7eed --- /dev/null +++ b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/core/main.cfg
@@ -0,0 +1,19 @@ +# Format: //devtools/kokoro/config/proto/build.proto + +# Copyright 2019 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Deliberately blank as everything necessary is configured in common files, but +# file must still exist to match corresponding (Google internal) job +# configurations that trigger the builds.
diff --git a/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/core/presubmit.cfg b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/core/presubmit.cfg new file mode 100755 index 0000000..50a7eed --- /dev/null +++ b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/core/presubmit.cfg
@@ -0,0 +1,19 @@ +# Format: //devtools/kokoro/config/proto/build.proto + +# Copyright 2019 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Deliberately blank as everything necessary is configured in common files, but +# file must still exist to match corresponding (Google internal) job +# configurations that trigger the builds.
diff --git a/kokoro/gcp_ubuntu/bazel/integrations/build.sh b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/integrations/build.sh similarity index 100% rename from kokoro/gcp_ubuntu/bazel/integrations/build.sh rename to build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/integrations/build.sh
diff --git a/kokoro/gcp_ubuntu/bazel/integrations/build_kokoro.sh b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/integrations/build_kokoro.sh old mode 100644 new mode 100755 similarity index 73% rename from kokoro/gcp_ubuntu/bazel/integrations/build_kokoro.sh rename to build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/integrations/build_kokoro.sh index 1716266..c35d897 --- a/kokoro/gcp_ubuntu/bazel/integrations/build_kokoro.sh +++ b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/integrations/build_kokoro.sh
@@ -16,28 +16,28 @@ # Build and test IREE's integrations within the gcr.io/iree-oss/bazel-tensorflow # image using Kokoro. +# Requires the environment variables KOKORO_ROOT and KOKORO_ARTIFACTS_DIR, which +# are set by Kokoro. -set -e set -x +set -e +set -o pipefail # Print the UTC time when set -x is on export PS4='[$(date -u "+%T %Z")] ' -# Kokoro checks out the repository here. -WORKDIR="${KOKORO_ARTIFACTS_DIR?}/github/iree" +source "${KOKORO_ARTIFACTS_DIR?}/github/iree/build_tools/kokoro/gcp_ubuntu/docker_common.sh" -# Mount the checked out repository, make that the working directory and run the -# tests in the bazel-tensorflow image. -docker run \ - --volume "${WORKDIR?}:${WORKDIR?}" \ - --workdir="${WORKDIR?}" \ - --rm \ +# Sets DOCKER_RUN_ARGS +docker_setup + +docker run "${DOCKER_RUN_ARGS[@]?}" \ gcr.io/iree-oss/bazel-tensorflow:prod \ - kokoro/gcp_ubuntu/bazel/integrations/build.sh + build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/integrations/build.sh # Kokoro will rsync this entire directory back to the executor orchestrating the # build which takes forever and is totally useless. -sudo rm -rf "${KOKORO_ARTIFACTS_DIR?}"/* +rm -rf "${KOKORO_ARTIFACTS_DIR?}"/* # Print out artifacts dir contents after deleting them as a coherence check. ls -1a "${KOKORO_ARTIFACTS_DIR?}/"
diff --git a/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/integrations/common.cfg b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/integrations/common.cfg new file mode 100644 index 0000000..eb31e55 --- /dev/null +++ b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/integrations/common.cfg
@@ -0,0 +1,20 @@ +# Format: //devtools/kokoro/config/proto/build.proto + +# Copyright 2019 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Common configuration for Kokoro builds that run the integrations build with +# bazel on linux. + +build_file: "iree/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/integrations/build_kokoro.sh"
diff --git a/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/integrations/continuous.cfg b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/integrations/continuous.cfg new file mode 100644 index 0000000..50a7eed --- /dev/null +++ b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/integrations/continuous.cfg
@@ -0,0 +1,19 @@ +# Format: //devtools/kokoro/config/proto/build.proto + +# Copyright 2019 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Deliberately blank as everything necessary is configured in common files, but +# file must still exist to match corresponding (Google internal) job +# configurations that trigger the builds.
diff --git a/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/integrations/google.cfg b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/integrations/google.cfg new file mode 100644 index 0000000..50a7eed --- /dev/null +++ b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/integrations/google.cfg
@@ -0,0 +1,19 @@ +# Format: //devtools/kokoro/config/proto/build.proto + +# Copyright 2019 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Deliberately blank as everything necessary is configured in common files, but +# file must still exist to match corresponding (Google internal) job +# configurations that trigger the builds.
diff --git a/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/integrations/main.cfg b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/integrations/main.cfg new file mode 100644 index 0000000..50a7eed --- /dev/null +++ b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/integrations/main.cfg
@@ -0,0 +1,19 @@ +# Format: //devtools/kokoro/config/proto/build.proto + +# Copyright 2019 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Deliberately blank as everything necessary is configured in common files, but +# file must still exist to match corresponding (Google internal) job +# configurations that trigger the builds.
diff --git a/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/integrations/presubmit.cfg b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/integrations/presubmit.cfg new file mode 100644 index 0000000..50a7eed --- /dev/null +++ b/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/integrations/presubmit.cfg
@@ -0,0 +1,19 @@ +# Format: //devtools/kokoro/config/proto/build.proto + +# Copyright 2019 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Deliberately blank as everything necessary is configured in common files, but +# file must still exist to match corresponding (Google internal) job +# configurations that trigger the builds.
diff --git a/kokoro/gcp_ubuntu/cmake/android/arm64-v8a/build_kokoro.sh b/build_tools/kokoro/gcp_ubuntu/cmake/android/arm64-v8a/build_kokoro.sh old mode 100644 new mode 100755 similarity index 73% rename from kokoro/gcp_ubuntu/cmake/android/arm64-v8a/build_kokoro.sh rename to build_tools/kokoro/gcp_ubuntu/cmake/android/arm64-v8a/build_kokoro.sh index f912073..329e226 --- a/kokoro/gcp_ubuntu/cmake/android/arm64-v8a/build_kokoro.sh +++ b/build_tools/kokoro/gcp_ubuntu/cmake/android/arm64-v8a/build_kokoro.sh
@@ -16,28 +16,28 @@ # Cross-compile the project towards Android arm64-v8a with the # gcr.io/iree-oss/cmake-android image using Kokoro. +# Requires the environment variables KOKORO_ROOT and KOKORO_ARTIFACTS_DIR, which +# are set by Kokoro. -set -e set -x +set -e +set -o pipefail # Print the UTC time when set -x is on export PS4='[$(date -u "+%T %Z")] ' -# Kokoro checks out the repository here. -WORKDIR=${KOKORO_ARTIFACTS_DIR?}/github/iree +source "${KOKORO_ARTIFACTS_DIR?}/github/iree/build_tools/kokoro/gcp_ubuntu/docker_common.sh" -# Mount the checked out repository, make that the working directory and run the -# tests in the cmake-android image. -docker run \ - --volume "${WORKDIR?}:${WORKDIR?}" \ - --workdir="${WORKDIR?}" \ - --rm \ +# Sets DOCKER_RUN_ARGS +docker_setup + +docker run "${DOCKER_RUN_ARGS[@]?}" \ gcr.io/iree-oss/cmake-android:prod \ - kokoro/gcp_ubuntu/cmake/android/build.sh arm64-v8a + build_tools/kokoro/gcp_ubuntu/cmake/android/build.sh arm64-v8a # Kokoro will rsync this entire directory back to the executor orchestrating the # build which takes forever and is totally useless. -sudo rm -rf "${KOKORO_ARTIFACTS_DIR?}"/* +rm -rf "${KOKORO_ARTIFACTS_DIR?}"/* # Print out artifacts dir contents after deleting them as a coherence check. ls -1a "${KOKORO_ARTIFACTS_DIR?}/"
diff --git a/build_tools/kokoro/gcp_ubuntu/cmake/android/arm64-v8a/common.cfg b/build_tools/kokoro/gcp_ubuntu/cmake/android/arm64-v8a/common.cfg new file mode 100644 index 0000000..1376e08 --- /dev/null +++ b/build_tools/kokoro/gcp_ubuntu/cmake/android/arm64-v8a/common.cfg
@@ -0,0 +1,20 @@ +# Format: //devtools/kokoro/config/proto/build.proto + +# Copyright 2020 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Common configuration for Kokoro builds that cross-compile IREE towards +# Android arm64-v8a using CMake. + +build_file: "iree/build_tools/kokoro/gcp_ubuntu/cmake/android/arm64-v8a/build_kokoro.sh"
diff --git a/build_tools/kokoro/gcp_ubuntu/cmake/android/arm64-v8a/continuous.cfg b/build_tools/kokoro/gcp_ubuntu/cmake/android/arm64-v8a/continuous.cfg new file mode 100644 index 0000000..e4cc270 --- /dev/null +++ b/build_tools/kokoro/gcp_ubuntu/cmake/android/arm64-v8a/continuous.cfg
@@ -0,0 +1,19 @@ +# Format: //devtools/kokoro/config/proto/build.proto + +# Copyright 2020 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Deliberately blank as everything necessary is configured in common files, but +# file must still exist to match corresponding (Google internal) job +# configurations that trigger the builds.
diff --git a/build_tools/kokoro/gcp_ubuntu/cmake/android/arm64-v8a/google.cfg b/build_tools/kokoro/gcp_ubuntu/cmake/android/arm64-v8a/google.cfg new file mode 100644 index 0000000..e4cc270 --- /dev/null +++ b/build_tools/kokoro/gcp_ubuntu/cmake/android/arm64-v8a/google.cfg
@@ -0,0 +1,19 @@ +# Format: //devtools/kokoro/config/proto/build.proto + +# Copyright 2020 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Deliberately blank as everything necessary is configured in common files, but +# file must still exist to match corresponding (Google internal) job +# configurations that trigger the builds.
diff --git a/build_tools/kokoro/gcp_ubuntu/cmake/android/arm64-v8a/main.cfg b/build_tools/kokoro/gcp_ubuntu/cmake/android/arm64-v8a/main.cfg new file mode 100644 index 0000000..e4cc270 --- /dev/null +++ b/build_tools/kokoro/gcp_ubuntu/cmake/android/arm64-v8a/main.cfg
@@ -0,0 +1,19 @@ +# Format: //devtools/kokoro/config/proto/build.proto + +# Copyright 2020 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Deliberately blank as everything necessary is configured in common files, but +# file must still exist to match corresponding (Google internal) job +# configurations that trigger the builds.
diff --git a/build_tools/kokoro/gcp_ubuntu/cmake/android/arm64-v8a/presubmit.cfg b/build_tools/kokoro/gcp_ubuntu/cmake/android/arm64-v8a/presubmit.cfg new file mode 100644 index 0000000..e4cc270 --- /dev/null +++ b/build_tools/kokoro/gcp_ubuntu/cmake/android/arm64-v8a/presubmit.cfg
@@ -0,0 +1,19 @@ +# Format: //devtools/kokoro/config/proto/build.proto + +# Copyright 2020 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Deliberately blank as everything necessary is configured in common files, but +# file must still exist to match corresponding (Google internal) job +# configurations that trigger the builds.
diff --git a/kokoro/gcp_ubuntu/cmake/android/build.sh b/build_tools/kokoro/gcp_ubuntu/cmake/android/build.sh similarity index 100% rename from kokoro/gcp_ubuntu/cmake/android/build.sh rename to build_tools/kokoro/gcp_ubuntu/cmake/android/build.sh
diff --git a/kokoro/gcp_ubuntu/cmake/linux/x86-turing/build.sh b/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86-turing/build.sh similarity index 87% rename from kokoro/gcp_ubuntu/cmake/linux/x86-turing/build.sh rename to build_tools/kokoro/gcp_ubuntu/cmake/linux/x86-turing/build.sh index a56c847..f44aa82 100755 --- a/kokoro/gcp_ubuntu/cmake/linux/x86-turing/build.sh +++ b/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86-turing/build.sh
@@ -30,8 +30,8 @@ python3 --version # Print Vulkan related information: SDK version and GPU ICD version -vulkaninfo 2>/dev/null | grep "Vulkan Instance" -vulkaninfo 2>/dev/null | grep -A7 "VkPhysicalDeviceProperties" +vulkaninfo 2>/dev/null | grep "Vulkan Instance" || echo "Vulkan Instance not found!" +vulkaninfo 2>/dev/null | grep -A7 "VkPhysicalDeviceProperties" || echo "VkPhysicalDeviceProperties not found!" echo "Initializing submodules" ./scripts/git/submodule_versions.py init
diff --git a/kokoro/gcp_ubuntu/cmake/linux/x86-turing/build_kokoro.sh b/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86-turing/build_kokoro.sh similarity index 64% rename from kokoro/gcp_ubuntu/cmake/linux/x86-turing/build_kokoro.sh rename to build_tools/kokoro/gcp_ubuntu/cmake/linux/x86-turing/build_kokoro.sh index 051c0d0..0b2364a 100755 --- a/kokoro/gcp_ubuntu/cmake/linux/x86-turing/build_kokoro.sh +++ b/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86-turing/build_kokoro.sh
@@ -14,32 +14,32 @@ # See the License for the specific language governing permissions and # limitations under the License. -# Build and test the project within the gcr.io/iree-oss/cmake using Kokoro. +# Build and test the project within the gcr.io/iree-oss/cmake image using +# Kokoro. +# Requires the environment variables KOKORO_ROOT and KOKORO_ARTIFACTS_DIR, which +# are set by Kokoro. -set -e set -x +set -e +set -o pipefail # Print the UTC time when set -x is on export PS4='[$(date -u "+%T %Z")] ' -# Kokoro checks out the repository here. -WORKDIR=${KOKORO_ARTIFACTS_DIR?}/github/iree +source "${KOKORO_ARTIFACTS_DIR?}/github/iree/build_tools/kokoro/gcp_ubuntu/docker_common.sh" -# Mount the checked out repository, make that the working directory and run the -# tests in the cmake image. -docker run \ - --volume "${WORKDIR?}:${WORKDIR?}" \ - --workdir="${WORKDIR?}" \ - --rm \ +# Sets DOCKER_RUN_ARGS +docker_setup + +docker run "${DOCKER_RUN_ARGS[@]?}" \ --env IREE_VULKAN_DISABLE=0 \ --gpus all \ gcr.io/iree-oss/cmake-nvidia:prod \ - kokoro/gcp_ubuntu/cmake/linux/x86-turing/build.sh + build_tools/kokoro/gcp_ubuntu/cmake/linux/x86-turing/build.sh # Kokoro will rsync this entire directory back to the executor orchestrating the # build which takes forever and is totally useless. -# TODO: enable this after making it work -#sudo rm -rf "${KOKORO_ARTIFACTS_DIR?}"/* +rm -rf "${KOKORO_ARTIFACTS_DIR?}"/* # Print out artifacts dir contents after deleting them as a coherence check. -#ls -1a "${KOKORO_ARTIFACTS_DIR?}/" +ls -1a "${KOKORO_ARTIFACTS_DIR?}/"
diff --git a/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86-turing/common.cfg b/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86-turing/common.cfg new file mode 100644 index 0000000..bdb9163 --- /dev/null +++ b/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86-turing/common.cfg
@@ -0,0 +1,19 @@ +# Format: //devtools/kokoro/config/proto/build.proto + +# Copyright 2020 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Common configuration for Kokoro builds that run cmake on linux. + +build_file: "iree/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86-turing/build_kokoro.sh"
diff --git a/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86-turing/continuous.cfg b/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86-turing/continuous.cfg new file mode 100644 index 0000000..e4cc270 --- /dev/null +++ b/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86-turing/continuous.cfg
@@ -0,0 +1,19 @@ +# Format: //devtools/kokoro/config/proto/build.proto + +# Copyright 2020 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Deliberately blank as everything necessary is configured in common files, but +# file must still exist to match corresponding (Google internal) job +# configurations that trigger the builds.
diff --git a/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86-turing/google.cfg b/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86-turing/google.cfg new file mode 100644 index 0000000..e4cc270 --- /dev/null +++ b/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86-turing/google.cfg
@@ -0,0 +1,19 @@ +# Format: //devtools/kokoro/config/proto/build.proto + +# Copyright 2020 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Deliberately blank as everything necessary is configured in common files, but +# file must still exist to match corresponding (Google internal) job +# configurations that trigger the builds.
diff --git a/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86-turing/main.cfg b/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86-turing/main.cfg new file mode 100644 index 0000000..e4cc270 --- /dev/null +++ b/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86-turing/main.cfg
@@ -0,0 +1,19 @@ +# Format: //devtools/kokoro/config/proto/build.proto + +# Copyright 2020 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Deliberately blank as everything necessary is configured in common files, but +# file must still exist to match corresponding (Google internal) job +# configurations that trigger the builds.
diff --git a/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86-turing/presubmit.cfg b/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86-turing/presubmit.cfg new file mode 100644 index 0000000..e4cc270 --- /dev/null +++ b/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86-turing/presubmit.cfg
@@ -0,0 +1,19 @@ +# Format: //devtools/kokoro/config/proto/build.proto + +# Copyright 2020 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Deliberately blank as everything necessary is configured in common files, but +# file must still exist to match corresponding (Google internal) job +# configurations that trigger the builds.
diff --git a/kokoro/gcp_ubuntu/cmake/build.sh b/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86/build.sh similarity index 100% rename from kokoro/gcp_ubuntu/cmake/build.sh rename to build_tools/kokoro/gcp_ubuntu/cmake/linux/x86/build.sh
diff --git a/kokoro/gcp_ubuntu/cmake/build_kokoro.sh b/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86/build_kokoro.sh similarity index 72% rename from kokoro/gcp_ubuntu/cmake/build_kokoro.sh rename to build_tools/kokoro/gcp_ubuntu/cmake/linux/x86/build_kokoro.sh index 8c35e60..ee459ff 100755 --- a/kokoro/gcp_ubuntu/cmake/build_kokoro.sh +++ b/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86/build_kokoro.sh
@@ -15,28 +15,28 @@ # limitations under the License. # Build and test the project within the gcr.io/iree-oss/cmake using Kokoro. +# Requires the environment variables KOKORO_ROOT and KOKORO_ARTIFACTS_DIR, which +# are set by Kokoro. -set -e set -x +set -e +set -o pipefail # Print the UTC time when set -x is on export PS4='[$(date -u "+%T %Z")] ' -# Kokoro checks out the repository here. -WORKDIR=${KOKORO_ARTIFACTS_DIR?}/github/iree +source "${KOKORO_ARTIFACTS_DIR?}/github/iree/build_tools/kokoro/gcp_ubuntu/docker_common.sh" -# Mount the checked out repository, make that the working directory and run the -# tests in the cmake image. -docker run \ - --volume "${WORKDIR?}:${WORKDIR?}" \ - --workdir="${WORKDIR?}" \ - --rm \ +# Sets DOCKER_RUN_ARGS +docker_setup + +docker run "${DOCKER_RUN_ARGS[@]?}" \ gcr.io/iree-oss/cmake:prod \ - kokoro/gcp_ubuntu/cmake/build.sh + build_tools/kokoro/gcp_ubuntu/cmake/linux/x86/build.sh # Kokoro will rsync this entire directory back to the executor orchestrating the # build which takes forever and is totally useless. -sudo rm -rf "${KOKORO_ARTIFACTS_DIR?}"/* +rm -rf "${KOKORO_ARTIFACTS_DIR?}"/* # Print out artifacts dir contents after deleting them as a coherence check. ls -1a "${KOKORO_ARTIFACTS_DIR?}/"
diff --git a/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86/common.cfg b/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86/common.cfg new file mode 100644 index 0000000..49e6865 --- /dev/null +++ b/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86/common.cfg
@@ -0,0 +1,19 @@ +# Format: //devtools/kokoro/config/proto/build.proto + +# Copyright 2020 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Common configuration for Kokoro builds that run cmake on linux. + +build_file: "iree/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86/build_kokoro.sh"
diff --git a/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86/continuous.cfg b/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86/continuous.cfg new file mode 100644 index 0000000..e4cc270 --- /dev/null +++ b/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86/continuous.cfg
@@ -0,0 +1,19 @@ +# Format: //devtools/kokoro/config/proto/build.proto + +# Copyright 2020 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Deliberately blank as everything necessary is configured in common files, but +# file must still exist to match corresponding (Google internal) job +# configurations that trigger the builds.
diff --git a/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86/google.cfg b/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86/google.cfg new file mode 100644 index 0000000..e4cc270 --- /dev/null +++ b/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86/google.cfg
@@ -0,0 +1,19 @@ +# Format: //devtools/kokoro/config/proto/build.proto + +# Copyright 2020 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Deliberately blank as everything necessary is configured in common files, but +# file must still exist to match corresponding (Google internal) job +# configurations that trigger the builds.
diff --git a/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86/main.cfg b/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86/main.cfg new file mode 100644 index 0000000..e4cc270 --- /dev/null +++ b/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86/main.cfg
@@ -0,0 +1,19 @@ +# Format: //devtools/kokoro/config/proto/build.proto + +# Copyright 2020 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Deliberately blank as everything necessary is configured in common files, but +# file must still exist to match corresponding (Google internal) job +# configurations that trigger the builds.
diff --git a/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86/presubmit.cfg b/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86/presubmit.cfg new file mode 100644 index 0000000..e4cc270 --- /dev/null +++ b/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86/presubmit.cfg
@@ -0,0 +1,19 @@ +# Format: //devtools/kokoro/config/proto/build.proto + +# Copyright 2020 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Deliberately blank as everything necessary is configured in common files, but +# file must still exist to match corresponding (Google internal) job +# configurations that trigger the builds.
diff --git a/build_tools/kokoro/gcp_ubuntu/docker_common.sh b/build_tools/kokoro/gcp_ubuntu/docker_common.sh new file mode 100644 index 0000000..81dc0b3 --- /dev/null +++ b/build_tools/kokoro/gcp_ubuntu/docker_common.sh
@@ -0,0 +1,64 @@ +# Copyright 2020 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Functions for setting up Docker containers to run on Kokoro + +# Sets up files and environment to enable running all our Kokoro docker scripts. +# In particular, does some shenanigans to enable running with the current user. +# Some of this setup is only strictly necessary for Bazel, but it doesn't hurt +# for anything else. +# Requires that KOKORO_ROOT and KOKORO_ARTIFACTS_DIR have been set +# Sets the environment variable DOCKER_RUN_ARGS to be used by subsequent +# `docker run` invocations. +function docker_setup() { + # Setup to run docker as the current user. + # Provide a place to mount files that would normally be under /etc/ + # We don't just mount the real /etc/passwd and /etc/group because Google + # Linux workstations do some interesting stuff with user/group permissions + # such that they don't contain the information about normal users and we + # want these scripts to be runnable locally for debugging. + local fake_etc_dir="${KOKORO_ROOT?}/fake_etc" + mkdir -p "${fake_etc_dir?}" + + local fake_group="${fake_etc_dir?}/group" + local fake_passwd="${fake_etc_dir?}/group" + + cp /etc/passwd "${fake_group?}" + cp /etc/group "${fake_passwd?}" + getent group "$(id -g)" >> "${fake_group?}" + getent passwd "$(id -u)" >> "${fake_passwd?}" + + + local workdir="${KOKORO_ARTIFACTS_DIR?}/github/iree" + + DOCKER_RUN_ARGS=( + # Run as the current user and group + --user="$(id -u):$(id -g)" + # Make the source repository available + --volume="${workdir?}:${workdir?}" + --workdir="${workdir?}" + # Tell docker about the host users and groups. Bazel needs this + # information, but it also makes some other things more pleasant. + --volume="${fake_group?}:/etc/group:ro" + --volume="${fake_passwd?}:/etc/passwd:ro" + # Allow Bazel to write its special cache directories. This is the + # default path Bazel will write to. + --volume="${HOME?}/.cache/bazel:${HOME?}/.cache/bazel" + # Make gcloud credentials available. This isn't necessary when running + # in GCE but enables using this script locally with RBE. + --volume="${HOME?}/.config/gcloud:${HOME?}/.config/gcloud:ro" + # Delete the container after + --rm + ) +}
diff --git a/build_tools/kokoro/gcp_ubuntu/simulate_kokoro.sh b/build_tools/kokoro/gcp_ubuntu/simulate_kokoro.sh new file mode 100755 index 0000000..38c6782 --- /dev/null +++ b/build_tools/kokoro/gcp_ubuntu/simulate_kokoro.sh
@@ -0,0 +1,54 @@ +#!/bin/bash + +# Copyright 2020 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +# Simulates the behavior of Kokoro on a local machine. +# Usage: +# ./kokoro/gcp_ubuntu/simulate_kokoro.sh build_tools/kokoro/gcp_ubuntu/bazel/core/build.sh +# +# Just does the part of the Kokoro setup that we care about and invokes the +# given build script. +# An optional second parameter can be used to specify a different repo to clone +# from. Especially useful for cloning the current git repo. +# ./kokoro/gcp_ubuntu/simulate_kokoro.sh build_tools/kokoro/gcp_ubuntu/bazel/core/build.sh "$PWD/.git" + +set -x +set -e +set -o pipefail + +RELATIVE_KOKORO_BUILD_SCRIPT="${1?}" +REPO_TO_CLONE="${2:-git@github.com:google/iree.git}" + +# Set up the temporary Kokoro directories +export KOKORO_ROOT="$(mktemp --directory --tmpdir kokoro-root-XXXXXX)" +mkdir -p "${KOKORO_ROOT?}/src/github" +export KOKORO_ARTIFACTS_DIR="${KOKORO_ROOT?}/src" +cd "${KOKORO_ARTIFACTS_DIR?}/github" + +# Clone the repo +git clone "${REPO_TO_CLONE?}" + +# The build script is assumed to be relative to the iree repo root. +KOKORO_BUILD_SCRIPT="${KOKORO_ARTIFACTS_DIR?}/github/iree/${RELATIVE_KOKORO_BUILD_SCRIPT?}" +chmod +x "${KOKORO_BUILD_SCRIPT?}" + +# This is where Kokoro starts its execution. +cd "${KOKORO_ARTIFACTS_DIR?}" + +# Run the actual script. +"${KOKORO_BUILD_SCRIPT?}" + +# Clean up after ourselves. +rm -rf "${KOKORO_ROOT?}"
diff --git a/build_tools/manylinux_py_setup.py b/build_tools/manylinux_py_setup.py new file mode 100755 index 0000000..abdd11d --- /dev/null +++ b/build_tools/manylinux_py_setup.py
@@ -0,0 +1,86 @@ +#!/opt/python/cp38-cp38/bin/python3 +# Copyright 2020 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. +"""Generates CMake arguments to build all manylinux python versions. + +manylinux containers have all python version linked under /opt/python. +This script scrapes them to get configuration, install deps, etc. + +Usage: + Install dependencies: + manylinux_py_setup.py deps + Get CMake arguments to build (typically via $() expansion): + manylinux_py_setup.py args +""" + +import os +from pathlib import Path +import subprocess +import sys +import sysconfig + + +def _get_python_exes(): + PYTHON_PARENT_PATH = Path("/opt/python") + return PYTHON_PARENT_PATH.glob("*/bin/python") + + +def install_deps(): + for python_exe in _get_python_exes(): + args = [ + str(python_exe), + "-m", + "pip", + "install", + "absl-py", + "numpy", + ] + print("EXEC:", " ".join(args)) + subprocess.check_call(args) + + +def dump_current(identifier): + print("-DIREE_MULTIPY_{}_EXECUTABLE='{}'".format(identifier, sys.executable)) + print("-DIREE_MULTIPY_{}_INCLUDE_DIRS='{}'".format( + identifier, sysconfig.get_config_var("INCLUDEPY"))) + # TODO: Print LIBRARIES for Windows and OSX + print("-DIREE_MULTIPY_{}_EXTENSION='{}'".format( + identifier, sysconfig.get_config_var("EXT_SUFFIX"))) + + +def dump_all(): + versions_ids = [] + for python_exe in _get_python_exes(): + identifier = python_exe.parent.parent.name + versions_ids.append(identifier) + # Invoke ourselves with a different interpreter/args to dump config. + subprocess.check_call( + [str(python_exe), __file__, "_current_args", identifier]) + print("-DIREE_MULTIPY_VERSIONS='{}'".format(";".join(versions_ids))) + + +if __name__ == "__main__": + if len(sys.argv) < 2: + print("SYNTAX: mainlinux_py_setup.py {deps|args}") + sys.exit(1) + command = sys.argv[1] + if command == "args": + dump_all() + elif command == "_current_args": + dump_current(sys.argv[2]) + elif command == "deps": + install_deps() + else: + print("Unexpected command") + sys.exit(1)
diff --git a/build_tools/third_party/flatcc/BUILD.overlay b/build_tools/third_party/flatcc/BUILD.overlay new file mode 100644 index 0000000..3481995 --- /dev/null +++ b/build_tools/third_party/flatcc/BUILD.overlay
@@ -0,0 +1,99 @@ +# Copyright 2020 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +package(default_visibility = ["//visibility:public"]) + +# NOTE: we exclude JSON parsing/printing to avoid additional dependencies. +cc_library( + name = "runtime", + srcs = [ + "config/config.h", + "src/runtime/builder.c", + "src/runtime/emitter.c", + "src/runtime/refmap.c", + "src/runtime/verifier.c", + ], + hdrs = [ + "include/flatcc/flatcc_accessors.h", + "include/flatcc/flatcc_alloc.h", + "include/flatcc/flatcc_assert.h", + "include/flatcc/flatcc_builder.h", + "include/flatcc/flatcc_emitter.h", + "include/flatcc/flatcc_endian.h", + "include/flatcc/flatcc_epilogue.h", + "include/flatcc/flatcc_flatbuffers.h", + "include/flatcc/flatcc_identifier.h", + "include/flatcc/flatcc_iov.h", + "include/flatcc/flatcc_portable.h", + "include/flatcc/flatcc_prologue.h", + "include/flatcc/flatcc_refmap.h", + "include/flatcc/flatcc_rtconfig.h", + "include/flatcc/flatcc_types.h", + "include/flatcc/flatcc_unaligned.h", + "include/flatcc/flatcc_verifier.h", + "include/flatcc/reflection/flatbuffers_common_builder.h", + "include/flatcc/reflection/flatbuffers_common_reader.h", + ] + glob(["include/flatcc/portable/**/*.h"]), + copts = [ + "-Iexternal/com_github_dvidelabs_flatcc/config/", + "-Iexternal/com_github_dvidelabs_flatcc/include/", + ], + includes = [ + "include/", + ], + strip_include_prefix = "include", +) + +cc_library( + name = "compiler", + srcs = glob([ + "external/**/*.c", + "external/**/*.h", + "src/compiler/**/*.c", + "src/compiler/**/*.h", + ], exclude = [ + "external/lex/luthor.c", + "**/*_test.c", + ]), + hdrs = glob([ + "config/config.h", + "include/**/*.h", + ]), + textual_hdrs = [ + "external/lex/luthor.c", + ] + glob([ + "external/**/*.h", + ]), + copts = [ + "-Iexternal/com_github_dvidelabs_flatcc/config/", + "-Iexternal/com_github_dvidelabs_flatcc/external/", + "-Iexternal/com_github_dvidelabs_flatcc/include/", + ], +) + +cc_binary( + name = "flatcc", + srcs = [ + "src/cli/flatcc_cli.c", + ], + deps = [ + ":compiler", + ":runtime", + ], + copts = [ + "-Iexternal/com_github_dvidelabs_flatcc/config/", + "-Iexternal/com_github_dvidelabs_flatcc/external/", + "-Iexternal/com_github_dvidelabs_flatcc/include/", + ], +)
diff --git a/build_tools/third_party/flatcc/CMakeLists.txt b/build_tools/third_party/flatcc/CMakeLists.txt new file mode 100644 index 0000000..82b1cb2 --- /dev/null +++ b/build_tools/third_party/flatcc/CMakeLists.txt
@@ -0,0 +1,53 @@ +# Copyright 2020 Google LLC +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# https://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +set(FLATCC_ROOT "${IREE_ROOT_DIR}/third_party/flatcc/") + +# NOTE: we exclude JSON parsing/printing to avoid additional dependencies. +external_cc_library( + PACKAGE + flatcc + NAME + runtime + ROOT + ${FLATCC_ROOT} + INCLUDES + "${FLATCC_ROOT}/include" + SRCS + "src/runtime/builder.c" + "src/runtime/emitter.c" + "src/runtime/refmap.c" + "src/runtime/verifier.c" + HDRS + "include/flatcc/flatcc_accessors.h" + "include/flatcc/flatcc_alloc.h" + "include/flatcc/flatcc_assert.h" + "include/flatcc/flatcc_builder.h" + "include/flatcc/flatcc_emitter.h" + "include/flatcc/flatcc_endian.h" + "include/flatcc/flatcc_epilogue.h" + "include/flatcc/flatcc_flatbuffers.h" + "include/flatcc/flatcc_identifier.h" + "include/flatcc/flatcc_iov.h" + "include/flatcc/flatcc_portable.h" + "include/flatcc/flatcc_prologue.h" + "include/flatcc/flatcc_refmap.h" + "include/flatcc/flatcc_rtconfig.h" + "include/flatcc/flatcc_types.h" + "include/flatcc/flatcc_unaligned.h" + "include/flatcc/flatcc_verifier.h" + "include/flatcc/reflection/flatbuffers_common_builder.h" + "include/flatcc/reflection/flatbuffers_common_reader.h" + PUBLIC +)
diff --git a/build_tools/third_party/swiftshader/build_vk_swiftshader.sh b/build_tools/third_party/swiftshader/build_vk_swiftshader.sh old mode 100644 new mode 100755
diff --git a/colab/README.md b/colab/README.md index d35df2a..b0bfe52 100644 --- a/colab/README.md +++ b/colab/README.md
@@ -1,7 +1,7 @@ # Google Colaboratory (Colab) Notebooks To run these notebooks with a local runtime, refer to the -[Using Colab docs](../docs/using_colab.md). +[Using Colab docs](../docs/using_iree/using_colab.md). Hosted/remote runtimes are not yet supported.
diff --git a/docs/design_docs/codegen_passes.md b/docs/design_docs/codegen_passes.md new file mode 100644 index 0000000..e3fd371 --- /dev/null +++ b/docs/design_docs/codegen_passes.md
@@ -0,0 +1,640 @@ +# 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.target` 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.target "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<4x5xf32> + %1 = hal.interface.load.tensor @legacy_io::@arg1, + offset = %c0 : tensor<5x10xf32> + %2 = "mhlo.dot"(%0, %1) {precision_config = ["DEFAULT", "DEFAULT"]} : + (tensor<4x5xf32>, tensor<5x10xf32>) -> tensor<4x10xf32> + hal.interface.store.tensor %2, @legacy_io::@ret0, + offset = %c0 : tensor<4x10xf32> + return + } + hal.interface @legacy_io attributes {sym_visibility = "private"} { + hal.interface.binding @arg0, set=0, binding=0, + type="StorageBuffer", access="Read" + hal.interface.binding @arg1, set=0, binding=1, + type="StorageBuffer", access="Read" + hal.interface.binding @ret0, set=0, binding=2, + type="StorageBuffer", access="Write|Discard" + } + } +} +``` +<a name="snippet1"></a> +Snippet 1 : Dispatch region with matrix-matrix multiply operation. + +```mlir +hal.executable.target "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<10x5xf32> + %1 = hal.interface.load.tensor @legacy_io::@arg1, + offset = %c0 : tensor<10x5xf32> + %2 = hal.interface.load.tensor @legacy_io::@arg2, + offset = %c0 : tensor<10x5xf32> + %3 = "mhlo.add"(%0, %1) : + (tensor<10x5xf32>, tensor<10x5xf32>) -> tensor<10x5xf32> + %4 = "mhlo.multiply"(%3, %2) : + (tensor<10x5xf32>, tensor<10x5xf32>) -> tensor<10x5xf32> + hal.interface.store.tensor %4, @legacy_io::@ret0, + offset = %c0 : tensor<10x5xf32> + return + } + hal.interface @legacy_io attributes {sym_visibility = "private"} { + hal.interface.binding @arg0, set=0, binding=0, + type="StorageBuffer", access="Read" + hal.interface.binding @arg1, set=0, binding=1, + type="StorageBuffer", access="Read" + hal.interface.binding @arg2, set=0, binding=2, + type="StorageBuffer", access="Read" + hal.interface.binding @ret0, set=0, binding=3, + type="StorageBuffer", access="Write|Discard" + } + } +} +``` +<a name="snippet2"></a> +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: + + + +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 each of the `mhlo.add` and `mhlo.multiply` +operations are converted to `linalg.generic` operations on tensors. + +```mlir +#map0 = affine_map<(d0, d1) -> (d0, 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<10x5xf32>, tensor<10x5xf32> -> tensor<10x5xf32> +%4 = linalg.generic + {args_in = 2 : i64, args_out = 1 : i64, + indexing_maps = [#map0, #map0, #map0], + iterator_types = ["parallel", "parallel"]} %3, %2 { + ^bb0(%arg0: f32, %arg1: f32): // no predecessors + %5 = mulf %arg0, %arg1 : f32 + linalg.yield %5 : f32 + }: tensor<10x5xf32>, tensor<10x5xf32> -> tensor<10x5xf32> +``` +<a name="snippet3"></a> +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)> +%3 = linalg.generic + {args_in = 3 : i64, args_out = 1 : i64, + indexing_maps = [#map0, #map0, #map0, #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<?x5xf32>, tensor<?x5xf32>, tensor<?x5xf32> -> tensor<?x5xf32> +``` +<a name="snippet4"></a> +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<4x10xf32> + %c0 = constant 0 : index + %1 = iree.placeholder for "interface buffer" + {binding = @legacy_io::@arg0} : memref<4x5xf32> + %2 = iree.placeholder for "interface buffer" + {binding = @legacy_io::@arg1} : memref<5x10xf32> + %cst = constant 0.000000e+00 : f32 + linalg.matmul(%1, %2, %0) : + memref<4x5xf32>, memref<5x10xf32>, memref<4x10xf32> + return +} +``` +<a name="snippet5"></a> +Snippet 5 : Matrix-matrix multiply after conversion to +Linalg operation on `memref`s. + +```mlir +#map0 = affine_map<(d0, d1) -> (d0, d1)> +func @main_ex_dispatch() { + %0 = iree.placeholder for "interface buffer" + {binding = @legacy_io::@ret0} : memref<10x5xf32> + %c0 = constant 0 : index + %1 = iree.placeholder for "interface buffer" + {binding = @legacy_io::@arg0} : memref<10x5xf32> + %2 = iree.placeholder for "interface buffer" + {binding = @legacy_io::@arg1} : memref<10x5xf32> + %3 = iree.placeholder for "interface buffer" + {binding = @legacy_io::@arg2} : memref<10x5xf32> + linalg.generic + {args_in = 3 : i64, args_out = 1 : i64, + indexing_maps = [#map0, #map0, #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<10x5xf32>, memref<10x5xf32>, memref<10x5xf32>, memref<10x5xf32> + return +} +``` +<a name="snippet6"></a> +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. + + + +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 + %c0 = constant 0 : index + %c4 = constant 4 : index + %c10 = constant 10 : index + %0 = iree.placeholder for "interface buffer" + {binding = @legacy_io::@ret0} : memref<4x10xf32> + %1 = iree.placeholder for "interface buffer" + {binding = @legacy_io::@arg0} : memref<4x5xf32> + %2 = iree.placeholder for "interface buffer" + {binding = @legacy_io::@arg1} : memref<5x10xf32> + linalg.fill(%0, %cst) : memref<4x10xf32>, f32 + scf.parallel (%arg0, %arg1) = (%c0, %c0) to (%c4, %c10) step (%c8, %c8) { + scf.for %arg2 = %c0 to %c5 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 +} +``` +<a name="snippet7"></a> +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 ([LinalgMarker][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>}} { + %c96 = constant 96 : 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<96x96xf32> + %1 = iree.placeholder for "interface buffer" + {binding = @legacy_io::@arg1} : memref<96x96xf32> + %2 = iree.placeholder for "interface buffer" + {binding = @legacy_io::@ret0} : memref<96x96xf32> + scf.parallel (%arg0, %arg1) = (%c0, %c0) to (%c96, %c96) step (%c8, %c8) { + scf.for %arg2 = %c0 to %c96 step %c4 { + ... + %5 = subview %0[%arg0, %arg2]... + ... + %8 = subview %1[%arg2, %arg1]... + ... + %11 = subview %2[%arg0, %arg1]... + %12 = alloc(%c8, %c4) : memref<?x?xf32, 3> + %13 = subview %12[%c0, %c0]... + %14 = alloc(%c4, %c8) : memref<?x?xf32, 3> + %15 = subview %14[%c0, %c0]... + linalg.copy(%5, %13) {__internal_linalg_transform__ = "workgroup"} + : memref<?x?xf32, #map2>, memref<?x?xf32, #map2, 3> + spv.ControlBarrier "Workgroup", "Workgroup", "AcquireRelease" + linalg.copy(%8, %15) {__internal_linalg_transform__ = "workgroup"} + : memref<?x?xf32, #map2>, memref<?x?xf32, #map2, 3> + spv.ControlBarrier "Workgroup", "Workgroup", "AcquireRelease" + linalg.matmul {__internal_linalg_transform__ = "workgroup"} %13, %15, %11... + spv.ControlBarrier "Workgroup", "Workgroup", "AcquireRelease" + dealloc %12 : memref<?x?xf32, 3> + dealloc %14 : memref<?x?xf32, 3> + } + scf.yield + } + return +} +``` + +<a name="snippet8"></a> +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>}} { + %c5 = constant 5 : 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<4x10xf32> + %1 = iree.placeholder for "interface buffer" + {binding = @legacy_io::@arg0} : memref<4x5xf32> + %2 = iree.placeholder for "interface buffer" + {binding = @legacy_io::@arg1} : memref<5x10xf32> + %3 = "gpu.block_id"() {dimension = "x"} : () -> index + %4 = muli %3, %c8 : index + scf.for %arg0 = %c0 to %c5 step %c4 { + ... + %9 = subview %1[0, %arg0] + ... + %14 = subview %2[%arg0, %4] + %15 = subview %0[0, %4] + %16 = "gpu.thread_id"() {dimension = "x"} : () -> index + %17 = "gpu.thread_id"() {dimension = "y"} : () -> index + %18 = cmpi "slt", %17, %c4 : index + %19 = cmpi "slt", %16, %13 : index + %20 = and %18, %19 : i1 + scf.if %20 { + scf.for %arg1 = %c0 to %8 step %c1 { + %21 = load %9[%17, %arg1] : memref<4x?xf32, #map0> + %22 = load %14[%arg1, %16] : memref<?x?xf32, #map1> + %23 = load %15[%17, %16] : memref<4x?xf32, #map1> + %24 = mulf %21, %22 : f32 + %25 = addf %23, %24 : f32 + store %25, %15[%17, %16] : memref<4x?xf32, #map1> + } + } + } + return +} +``` +<a name="snippet9"></a> +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<10x5xf32> + %1 = iree.placeholder for "interface buffer" + {binding = @legacy_io::@arg0} : memref<10x5xf32> + %2 = iree.placeholder for "interface buffer" + {binding = @legacy_io::@arg1} : memref<10x5xf32> + %3 = iree.placeholder for "interface buffer" + {binding = @legacy_io::@arg2} : memref<10x5xf32> + %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<10x5xf32> + %13 = load %2[%10, %11] : memref<10x5xf32> + %14 = load %3[%10, %11] : memref<10x5xf32> + %15 = addf %12, %13 : f32 + %16 = mulf %15, %14 : f32 + store %16, %0[%10, %11] : memref<10x5xf32> + } + return +} +``` +<a name="snippet10"></a> +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/google/iree/blob/main/iree/compiler/Conversion/LinalgToSPIRV/ConvertToGPUPass.cpp +[ConvertToSPIRV]: https://github.com/google/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/google/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/dynamic_shapes.md b/docs/design_docs/dynamic_shapes.md similarity index 100% rename from docs/dynamic_shapes.md rename to docs/design_docs/dynamic_shapes.md
diff --git a/docs/function_abi.md b/docs/design_docs/function_abi.md similarity index 99% rename from docs/function_abi.md rename to docs/design_docs/function_abi.md index 797f2f6..71d91bb 100644 --- a/docs/function_abi.md +++ b/docs/design_docs/function_abi.md
@@ -1,4 +1,4 @@ -# Function signatures +# Function Signatures 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/design_docs/hlo_to_linalg.png b/docs/design_docs/hlo_to_linalg.png new file mode 100755 index 0000000..469ed26 --- /dev/null +++ b/docs/design_docs/hlo_to_linalg.png Binary files differ
diff --git a/docs/design_docs/linalg_to_spirv.png b/docs/design_docs/linalg_to_spirv.png new file mode 100755 index 0000000..fd6aee7 --- /dev/null +++ b/docs/design_docs/linalg_to_spirv.png Binary files differ
diff --git a/docs/simple_ir_walkthrough.md b/docs/design_docs/simple_ir_walkthrough.md similarity index 99% rename from docs/simple_ir_walkthrough.md rename to docs/design_docs/simple_ir_walkthrough.md index 1dfcb49..68f51a8 100644 --- a/docs/simple_ir_walkthrough.md +++ b/docs/design_docs/simple_ir_walkthrough.md
@@ -1,5 +1,7 @@ # Simple IR Walkthrough +Note that this doc is quite outdated. We expect to update it soon. + ## Overview This walks through the process of lowering TensorFlow python to an IREE module,
diff --git a/docs/roadmap_design.md b/docs/design_roadmap.md similarity index 100% rename from docs/roadmap_design.md rename to docs/design_roadmap.md
diff --git a/docs/benchmarking.md b/docs/developing_iree/benchmarking.md similarity index 100% rename from docs/benchmarking.md rename to docs/developing_iree/benchmarking.md
diff --git a/docs/contributor_tips.md b/docs/developing_iree/contributor_tips.md similarity index 89% rename from docs/contributor_tips.md rename to docs/developing_iree/contributor_tips.md index 05fb52e..03ddffc 100644 --- a/docs/contributor_tips.md +++ b/docs/developing_iree/contributor_tips.md
@@ -2,11 +2,11 @@ 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. +(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
diff --git a/docs/developer_overview.md b/docs/developing_iree/developer_overview.md similarity index 89% rename from docs/developer_overview.md rename to docs/developing_iree/developer_overview.md index 56fb5f8..ec900f5 100644 --- a/docs/developer_overview.md +++ b/docs/developing_iree/developer_overview.md
@@ -150,7 +150,7 @@ and executes it as a series of [googletest](https://github.com/google/googletest) tests. This is the test runner for the IREE -[check framework](https://github.com/google/iree/tree/main/docs/testing_guide.md#end-to-end-tests). +[check framework](https://github.com/google/iree/tree/main/docs/developing_iree/testing_guide.md#end-to-end-tests). ```shell $ bazel run iree/tools:iree-translate -- \ @@ -207,6 +207,28 @@ accept a number where 0, 1, 2, 3 stands for info, warning, error, and fatal error respectively. +#### 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 +scalar or a buffer. Scalars should be in the format `type=value` and buffers +should be in the format `[shape]xtype=[value]`. For example: + +``` +1x5xf32=1,-2,-3,4,-5 +1x5x3x1xf32=15,14,13,12,11,10,9,8,7,6,5,4,3,2,1 +``` + +#### `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 +multiple dispatch function. When the flag is on, IREE will insert trace points +before and after each dispatch function. The first trace op is for inputs, and +the second trace op is for outputs. There will be two events for one dispatch +function. + ### Useful Vulkan driver flags For IREE's Vulkan runtime driver, there are a few useful
diff --git a/docs/repository_management.md b/docs/developing_iree/repository_management.md similarity index 100% rename from docs/repository_management.md rename to docs/developing_iree/repository_management.md
diff --git a/docs/testing_guide.md b/docs/developing_iree/testing_guide.md similarity index 98% rename from docs/testing_guide.md rename to docs/developing_iree/testing_guide.md index 3be690d..5f57c9c 100644 --- a/docs/testing_guide.md +++ b/docs/developing_iree/testing_guide.md
@@ -30,7 +30,7 @@ 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](GetStarted/generic_vulkan_env_setup.md#useful-environment-variables) +[generic Vulkan setup](get_started/generic_vulkan_env_setup.md#useful-environment-variables) page for details regarding these variables. For Bazel, you can persist the configuration in `user.bazelrc` to save typing.
diff --git a/docs/GetStarted/cmake_options_and_variables.md b/docs/get_started/cmake_options_and_variables.md similarity index 91% rename from docs/GetStarted/cmake_options_and_variables.md rename to docs/get_started/cmake_options_and_variables.md index 8e121fc..3f3dbdd 100644 --- a/docs/GetStarted/cmake_options_and_variables.md +++ b/docs/get_started/cmake_options_and_variables.md
@@ -63,17 +63,21 @@ #### `IREE_HAL_DRIVERS_TO_BUILD`:STRING -*This does not have any effect at the moment, but will be supported in the -future!* Semicolon-separated list of HAL drivers to build, or `all` for building -all HAL drivers. Case-insensitive. Defaults to `all`. Example: +*Righ now this only affects whether tests are enabled when compiling for +Android; it will be fully supported in the future!* + +Semicolon-separated list of HAL drivers to build, or `all` for building all HAL +drivers. Case-insensitive. Defaults to `all`. Example: `-DIREE_HAL_DRIVERS_TO_BUILD="Vulkan;VMLA"`. #### `IREE_TARGET_BACKENDS_TO_BUILD`:STRING -*This does not have any effect at the moment, but will be supported in the -future!* Semicolon-separated list of HAL drivers to build, or `all` for building -all HAL drivers. Case-insensitive. Defaults to `all`. Example: -`-DIREE_HAL_DRIVERS_TO_BUILD="Vulkan_SPIRV;VMLA"`. +*Righ now this only affects whether tests are enabled when compiling for +Android; it will be fully supported in the future!* + +Semicolon-separated list of HAL drivers to build, or `all` for building all +compiler target backends. Case-insensitive. Defaults to `all`. Example: +`-DIREE_HAL_DRIVERS_TO_BUILD="Vulkan-SPIRV;VMLA"`. #### `IREE_ENABLE_LLD`:BOOL
diff --git a/docs/GetStarted/generic_vulkan_env_setup.md b/docs/get_started/generic_vulkan_env_setup.md similarity index 100% rename from docs/GetStarted/generic_vulkan_env_setup.md rename to docs/get_started/generic_vulkan_env_setup.md
diff --git a/docs/GetStarted/getting_started_android_cmake.md b/docs/get_started/getting_started_android_cmake.md similarity index 100% rename from docs/GetStarted/getting_started_android_cmake.md rename to docs/get_started/getting_started_android_cmake.md
diff --git a/docs/GetStarted/getting_started_linux_bazel.md b/docs/get_started/getting_started_linux_bazel.md similarity index 97% rename from docs/GetStarted/getting_started_linux_bazel.md rename to docs/get_started/getting_started_linux_bazel.md index 8d4a9dc..7f8b688 100644 --- a/docs/GetStarted/getting_started_linux_bazel.md +++ b/docs/get_started/getting_started_linux_bazel.md
@@ -123,7 +123,7 @@ ### Further Reading * For an introduction to IREE's project structure and developer tools, see - [Developer Overview](../developer_overview.md) + [Developer Overview](../developing_iree/developer_overview.md) * To target GPUs using Vulkan, see [Getting Started on Linux with Vulkan](getting_started_linux_vulkan.md) * To use IREE's Python bindings, see
diff --git a/docs/GetStarted/getting_started_linux_cmake.md b/docs/get_started/getting_started_linux_cmake.md similarity index 97% rename from docs/GetStarted/getting_started_linux_cmake.md rename to docs/get_started/getting_started_linux_cmake.md index 70da146..127a80c 100644 --- a/docs/GetStarted/getting_started_linux_cmake.md +++ b/docs/get_started/getting_started_linux_cmake.md
@@ -110,7 +110,7 @@ ### Further Reading * For an introduction to IREE's project structure and developer tools, see - [Developer Overview](../developer_overview.md) + [Developer Overview](../developing_iree/developer_overview.md) * To target GPUs using Vulkan, see [Getting Started on Linux with Vulkan](getting_started_linux_vulkan.md) * To use IREE's Python bindings, see
diff --git a/docs/GetStarted/getting_started_linux_vulkan.md b/docs/get_started/getting_started_linux_vulkan.md similarity index 100% rename from docs/GetStarted/getting_started_linux_vulkan.md rename to docs/get_started/getting_started_linux_vulkan.md
diff --git a/docs/GetStarted/getting_started_macos_bazel.md b/docs/get_started/getting_started_macos_bazel.md similarity index 95% rename from docs/GetStarted/getting_started_macos_bazel.md rename to docs/get_started/getting_started_macos_bazel.md index 2b285eb..3fc3dcb 100644 --- a/docs/GetStarted/getting_started_macos_bazel.md +++ b/docs/get_started/getting_started_macos_bazel.md
@@ -126,8 +126,8 @@ ### Further Reading * For an introduction to IREE's project structure and developer tools, see - [Developer Overview](../developer_overview.md) <!-- TODO: Link to macOS - versions of these guides once they are developed. + [Developer Overview](../developing_iree/developer_overview.md) <!-- TODO: + Link to macOS versions of these guides once they are developed. * To target GPUs using Vulkan, see [Getting Started on Linux with Vulkan](getting_started_linux_vulkan.md) * To use IREE's Python bindings, see
diff --git a/docs/GetStarted/getting_started_macos_cmake.md b/docs/get_started/getting_started_macos_cmake.md similarity index 95% rename from docs/GetStarted/getting_started_macos_cmake.md rename to docs/get_started/getting_started_macos_cmake.md index 51ef0ab..7b916cd 100644 --- a/docs/GetStarted/getting_started_macos_cmake.md +++ b/docs/get_started/getting_started_macos_cmake.md
@@ -110,8 +110,8 @@ ### Further Reading * For an introduction to IREE's project structure and developer tools, see - [Developer Overview](../developer_overview.md) <!-- TODO: Link to macOS - versions of these guides once they are developed. + [Developer Overview](../developing_iree/developer_overview.md) <!-- TODO: + Link to macOS versions of these guides once they are developed. * To target GPUs using Vulkan, see [Getting Started on Linux with Vulkan](getting_started_linux_vulkan.md) * To use IREE's Python bindings, see
diff --git a/docs/GetStarted/getting_started_python.md b/docs/get_started/getting_started_python.md similarity index 100% rename from docs/GetStarted/getting_started_python.md rename to docs/get_started/getting_started_python.md
diff --git a/docs/GetStarted/getting_started_windows_bazel.md b/docs/get_started/getting_started_windows_bazel.md similarity index 97% rename from docs/GetStarted/getting_started_windows_bazel.md rename to docs/get_started/getting_started_windows_bazel.md index d3f01c3..8cf0f87 100644 --- a/docs/GetStarted/getting_started_windows_bazel.md +++ b/docs/get_started/getting_started_windows_bazel.md
@@ -118,7 +118,7 @@ ### Further Reading * For an introduction to IREE's project structure and developer tools, see - [Developer Overview](../developer_overview.md) + [Developer Overview](../developing_iree/developer_overview.md) * To target GPUs using Vulkan, see [Getting Started on Windows with Vulkan](getting_started_windows_vulkan.md) * To use IREE's Python bindings, see
diff --git a/docs/GetStarted/getting_started_windows_cmake.md b/docs/get_started/getting_started_windows_cmake.md similarity index 97% rename from docs/GetStarted/getting_started_windows_cmake.md rename to docs/get_started/getting_started_windows_cmake.md index dcab418..da5218f 100644 --- a/docs/GetStarted/getting_started_windows_cmake.md +++ b/docs/get_started/getting_started_windows_cmake.md
@@ -107,7 +107,7 @@ ### Further Reading * For an introduction to IREE's project structure and developer tools, see - [Developer Overview](../developer_overview.md) + [Developer Overview](../developing_iree/developer_overview.md) * To target GPUs using Vulkan, see [Getting Started on Windows with Vulkan](getting_started_windows_vulkan.md) * To use IREE's Python bindings, see
diff --git a/docs/GetStarted/getting_started_windows_vulkan.md b/docs/get_started/getting_started_windows_vulkan.md similarity index 100% rename from docs/GetStarted/getting_started_windows_vulkan.md rename to docs/get_started/getting_started_windows_vulkan.md
diff --git a/docs/IREE-Architecture.svg b/docs/iree_architecture.svg similarity index 100% rename from docs/IREE-Architecture.svg rename to docs/iree_architecture.svg
diff --git a/docs/roadmap.md b/docs/milestones.md similarity index 94% rename from docs/roadmap.md rename to docs/milestones.md index b979de7..322fa7e 100644 --- a/docs/roadmap.md +++ b/docs/milestones.md
@@ -1,11 +1,11 @@ -# IREE Roadmap +# IREE Milestones ## Design Though many of the core dialects are now in place enough for correctness testing a large majority of the features we are most excited to demonstrate are still TODO and will be coming over the next few quarters. You can find a highlighted -set of coming features in the [design roadmap](roadmap_design.md). +set of coming features in the [design roadmap](design_roadmap.md). ## Spring/Summer 2020 Focus Areas @@ -37,7 +37,7 @@ ### HAL: Marl CPU Scheduling We want to plug in [marl](https://github.com/google/marl) to provide -[CPU-side work scheduling](roadmap_design.md#gpu-like-cpu-scheduling) that +[CPU-side work scheduling](design_roadmap.md#gpu-like-cpu-scheduling) that matches GPU semantics. This will enable improved CPU utilization and allow us to verify the approach with benchmarks.
diff --git a/docs/mnist_example.md b/docs/mnist_example.md deleted file mode 100644 index 93cd23a..0000000 --- a/docs/mnist_example.md +++ /dev/null
@@ -1,254 +0,0 @@ -# MNIST IR Example - -This shows the MNIST MLP model as it is compiled from Keras, lowered to XLA HLO, -and then lowered to an IREE module with SPIR-V. Several steps are omitted for -brevity. - -## TensorFlow Keras Model - -```python -def simple_mnist_model(input_shape): - """Creates a simple (multi-layer perceptron) MNIST model.""" - model = tf.keras.models.Sequential() - # Flatten to a 1d array (e.g. 28x28 -> 784) - model.add(tf.keras.layers.Flatten(input_shape=input_shape)) - # Fully-connected neural layer with 128 neurons, RELU activation - model.add(tf.keras.layers.Dense(128, activation='relu')) - # Fully-connected neural layer returning probability scores for each class - model.add(tf.keras.layers.Dense(10, activation='softmax')) - return model -``` - -## XLA HLO - -**NOTE**: this uses placeholder weights to keep the page from being a few -thousand lines of floats. - -```mlir -module { - func @main(%arg0: tensor<1x28x28x1xf32>) -> tuple<tensor<1x10xf32>> - attributes {iree.module.export} { - %cst = constant {name = "constant.9"} dense<0.5> : tensor<f32> - %0 = "mhlo.broadcast_in_dim"(%cst) {name = "broadcast.10"} : (tensor<f32>) -> tensor<1x128xf32> - %1 = "mhlo.copy"(%arg0) {name = "copy.1"} : (tensor<1x28x28x1xf32>) -> tensor<1x28x28x1xf32> - %2 = "mhlo.reshape"(%1) {name = "reshape.2"} : (tensor<1x28x28x1xf32>) -> tensor<1x28x28x1xf32> - %3 = "mhlo.reshape"(%2) {name = "reshape.3"} : (tensor<1x28x28x1xf32>) -> tensor<1x784xf32> - %cst_0 = constant {name = "constant.4"} dense<0.5> : tensor<784x128xf32> - %4 = "mhlo.dot"(%3, %cst_0) {name = "dot.5", precision_config = ["DEFAULT", "DEFAULT"]} : (tensor<1x784xf32>, tensor<784x128xf32>) -> tensor<1x128xf32> - %cst_1 = constant {name = "constant.6"} dense<0.5> : tensor<128xf32> - %5 = "mhlo.broadcast_in_dim"(%cst_1) {broadcast_dimensions = dense<1> : tensor<1xi64>, name = "broadcast.7"} : (tensor<128xf32>) -> tensor<1x128xf32> - %6 = "mhlo.add"(%4, %5) {name = "add.8"} : (tensor<1x128xf32>, tensor<1x128xf32>) -> tensor<1x128xf32> - %7 = "mhlo.maximum"(%0, %6) {name = "maximum.11"} : (tensor<1x128xf32>, tensor<1x128xf32>) -> tensor<1x128xf32> - %cst_2 = constant {name = "constant.12"} dense<0.5> : tensor<128x10xf32> - %8 = "mhlo.dot"(%7, %cst_2) {name = "dot.13", precision_config = ["DEFAULT", "DEFAULT"]} : (tensor<1x128xf32>, tensor<128x10xf32>) -> tensor<1x10xf32> - %cst_3 = constant {name = "constant.14"} dense<0.5> : tensor<10xf32> - %9 = "mhlo.broadcast_in_dim"(%cst_3) {broadcast_dimensions = dense<1> : tensor<1xi64>, name = "broadcast.15"} : (tensor<10xf32>) -> tensor<1x10xf32> - %10 = "mhlo.add"(%8, %9) {name = "add.16"} : (tensor<1x10xf32>, tensor<1x10xf32>) -> tensor<1x10xf32> - %cst_4 = constant {name = "constant.17"} dense<0xFF800000> : tensor<f32> - %11 = "mhlo.reduce"(%10, %cst_4) ( { - ^bb0(%arg1: tensor<f32>, %arg2: tensor<f32>): // no predecessors - %20 = "mhlo.maximum"(%arg1, %arg2) {name = "maximum.21"} : (tensor<f32>, tensor<f32>) -> tensor<f32> - "mhlo.return"(%20) : (tensor<f32>) -> () - }) {dimensions = dense<1> : tensor<1xi64>} : (tensor<1x10xf32>, tensor<f32>) -> tensor<1xf32> - %12 = "mhlo.broadcast_in_dim"(%11) {broadcast_dimensions = dense<0> : tensor<1xi64>, name = "broadcast.23"} : (tensor<1xf32>) -> tensor<1x10xf32> - %13 = "mhlo.subtract"(%10, %12) {name = "subtract.24"} : (tensor<1x10xf32>, tensor<1x10xf32>) -> tensor<1x10xf32> - %14 = "mhlo.exponential"(%13) {name = "exponential.25"} : (tensor<1x10xf32>) -> tensor<1x10xf32> - %cst_5 = constant {name = "constant.27"} dense<0.5> : tensor<f32> - %15 = "mhlo.reduce"(%14, %cst_5) ( { - ^bb0(%arg3: tensor<f32>, %arg4: tensor<f32>): // no predecessors - %21 = "mhlo.add"(%arg3, %arg4) {name = "add.31"} : (tensor<f32>, tensor<f32>) -> tensor<f32> - "mhlo.return"(%21) : (tensor<f32>) -> () - }) {dimensions = dense<1> : tensor<1xi64>} : (tensor<1x10xf32>, tensor<f32>) -> tensor<1xf32> - %16 = "mhlo.broadcast_in_dim"(%15) {broadcast_dimensions = dense<0> : tensor<1xi64>, name = "broadcast.34"} : (tensor<1xf32>) -> tensor<1x10xf32> - %17 = "mhlo.divide"(%14, %16) {name = "divide.35"} : (tensor<1x10xf32>, tensor<1x10xf32>) -> tensor<1x10xf32> - %18 = "mhlo.reshape"(%17) {name = "reshape.36"} : (tensor<1x10xf32>) -> tensor<1x10xf32> - %19 = "mhlo.tuple"(%18) {name = "tuple.37"} : (tensor<1x10xf32>) -> tuple<tensor<1x10xf32>> - return %19 : tuple<tensor<1x10xf32>> - } -} -``` - -## IREE IR (pre-backend lowering) - -Here's the lowered, outlined, and compiler-annotated version of the above in the -IREE sequencer dialect. - -```mlir -module { - iree.multi_arch_executable @main_ex_dispatch_0[0]() { - iree.executable[0](Unspecified) { - module { - func @main_entry_dispatch_0(%arg0: memref<1x28x28x1xf32>, %arg1: memref<1x784xf32>) - attributes {iree.executable.export, iree.executable.workload = dense<[784, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { - %0 = iree.load_input(%arg0 : memref<1x28x28x1xf32>) : tensor<1x28x28x1xf32> - %1 = "mhlo.copy"(%0) {name = "copy.1"} : (tensor<1x28x28x1xf32>) -> tensor<1x28x28x1xf32> - %2 = "mhlo.reshape"(%1) {name = "reshape.3"} : (tensor<1x28x28x1xf32>) -> tensor<1x784xf32> - iree.store_output(%2 : tensor<1x784xf32>, %arg1 : memref<1x784xf32>) - iree.return - } - } - } - } - iree.multi_arch_executable @main_ex_dispatch_1[1]() { - iree.executable[1](Unspecified) { - module { - func @main_entry_dispatch_1(%arg0: memref<1x784xf32>, %arg1: memref<784x128xf32>, %arg2: memref<1x128xf32>) - attributes {iree.executable.export, iree.executable.workload = dense<[128, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { - %0 = iree.load_input(%arg0 : memref<1x784xf32>) : tensor<1x784xf32> - %1 = iree.load_input(%arg1 : memref<784x128xf32>) : tensor<784x128xf32> - %2 = "mhlo.dot"(%0, %1) {name = "dot.5", precision_config = ["DEFAULT", "DEFAULT"]} : (tensor<1x784xf32>, tensor<784x128xf32>) -> tensor<1x128xf32> - iree.store_output(%2 : tensor<1x128xf32>, %arg2 : memref<1x128xf32>) - iree.return - } - } - } - } - iree.multi_arch_executable @main_ex_dispatch_2[2]() { - iree.executable[2](Unspecified) { - module { - func @main_entry_dispatch_2(%arg0: memref<1x128xf32>, %arg1: memref<1x128xf32>) - attributes {iree.executable.export, iree.executable.workload = dense<[128, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { - %0 = iree.load_input(%arg0 : memref<1x128xf32>) : tensor<1x128xf32> - %cst = constant dense<5.000000e-01> : tensor<128xf32> - %cst_0 = constant dense<5.000000e-01> : tensor<f32> - %1 = "mhlo.broadcast_in_dim"(%cst_0) {name = "broadcast.10"} : (tensor<f32>) -> tensor<1x128xf32> - %2 = "mhlo.broadcast_in_dim"(%cst) {broadcast_dimensions = dense<1> : tensor<1xi64>, name = "broadcast.7"} : (tensor<128xf32>) -> tensor<1x128xf32> - %3 = addf %0, %2 : tensor<1x128xf32> - %4 = mhlo.maximum %1, %3 {name = "maximum.11"} : tensor<1x128xf32> - iree.store_output(%4 : tensor<1x128xf32>, %arg1 : memref<1x128xf32>) - iree.return - } - } - } - } - iree.multi_arch_executable @main_ex_dispatch_3[3]() { - iree.executable[3](Unspecified) { - module { - func @main_entry_dispatch_3(%arg0: memref<1x128xf32>, %arg1: memref<128x10xf32>, %arg2: memref<1x10xf32>) - attributes {iree.executable.export, iree.executable.workload = dense<[10, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { - %0 = iree.load_input(%arg0 : memref<1x128xf32>) : tensor<1x128xf32> - %1 = iree.load_input(%arg1 : memref<128x10xf32>) : tensor<128x10xf32> - %2 = "mhlo.dot"(%0, %1) {name = "dot.13", precision_config = ["DEFAULT", "DEFAULT"]} : (tensor<1x128xf32>, tensor<128x10xf32>) -> tensor<1x10xf32> - iree.store_output(%2 : tensor<1x10xf32>, %arg2 : memref<1x10xf32>) - iree.return - } - } - } - } - iree.multi_arch_executable @main_ex_dispatch_4[4]() { - iree.executable[4](Unspecified) { - module { - func @main_entry_dispatch_4(%arg0: memref<1x10xf32>, %arg1: memref<1x10xf32>) - attributes {iree.executable.export, iree.executable.workload = dense<[10, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { - %0 = iree.load_input(%arg0 : memref<1x10xf32>) : tensor<1x10xf32> - %cst = constant dense<5.000000e-01> : tensor<10xf32> - %1 = "mhlo.broadcast_in_dim"(%cst) {broadcast_dimensions = dense<1> : tensor<1xi64>, name = "broadcast.15"} : (tensor<10xf32>) -> tensor<1x10xf32> - %2 = addf %0, %1 : tensor<1x10xf32> - iree.store_output(%2 : tensor<1x10xf32>, %arg1 : memref<1x10xf32>) - iree.return - } - } - } - } - iree.multi_arch_executable @main_ex_dispatch_5[5]() { - iree.executable[5](Unspecified) { - module { - func @main_entry_dispatch_5(%arg0: memref<1x10xf32>, %arg1: memref<1xf32>) - attributes {iree.executable.export, iree.executable.workload = dense<1> : tensor<3xi32>, iree.ordinal = 0 : i32} { - %0 = iree.load_input(%arg0 : memref<1x10xf32>) : tensor<1x10xf32> - %cst = constant dense<0xFF800000> : tensor<f32> - %1 = "mhlo.reduce"(%0, %cst) ( { - ^bb0(%arg2: tensor<f32>, %arg3: tensor<f32>): // no predecessors - %2 = mhlo.maximum %arg2, %arg3 {name = "maximum.21"} : tensor<f32> - "mhlo.return"(%2) : (tensor<f32>) -> () - }) {dimensions = dense<1> : tensor<1xi64>} : (tensor<1x10xf32>, tensor<f32>) -> tensor<1xf32> - iree.store_output(%1 : tensor<1xf32>, %arg1 : memref<1xf32>) - iree.return - } - } - } - } - iree.multi_arch_executable @main_ex_dispatch_6[6]() { - iree.executable[6](Unspecified) { - module { - func @main_entry_dispatch_6(%arg0: memref<1x10xf32>, %arg1: memref<1xf32>, %arg2: memref<1x10xf32>) - attributes {iree.executable.export, iree.executable.workload = dense<[10, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { - %0 = iree.load_input(%arg0 : memref<1x10xf32>) : tensor<1x10xf32> - %1 = iree.load_input(%arg1 : memref<1xf32>) : tensor<1xf32> - %2 = "mhlo.broadcast_in_dim"(%1) {broadcast_dimensions = dense<0> : tensor<1xi64>, name = "broadcast.23"} : (tensor<1xf32>) -> tensor<1x10xf32> - %3 = subf %0, %2 : tensor<1x10xf32> - %4 = "mhlo.exponential"(%3) {name = "exponential.25"} : (tensor<1x10xf32>) -> tensor<1x10xf32> - iree.store_output(%4 : tensor<1x10xf32>, %arg2 : memref<1x10xf32>) - iree.return - } - } - } - } - iree.multi_arch_executable @main_ex_dispatch_7[7]() { - iree.executable[7](Unspecified) { - module { - func @main_entry_dispatch_7(%arg0: memref<1x10xf32>, %arg1: memref<1xf32>) - attributes {iree.executable.export, iree.executable.workload = dense<1> : tensor<3xi32>, iree.ordinal = 0 : i32} { - %0 = iree.load_input(%arg0 : memref<1x10xf32>) : tensor<1x10xf32> - %cst = constant dense<5.000000e-01> : tensor<f32> - %1 = "mhlo.reduce"(%0, %cst) ( { - ^bb0(%arg2: tensor<f32>, %arg3: tensor<f32>): // no predecessors - %2 = addf %arg2, %arg3 : tensor<f32> - "mhlo.return"(%2) : (tensor<f32>) -> () - }) {dimensions = dense<1> : tensor<1xi64>} : (tensor<1x10xf32>, tensor<f32>) -> tensor<1xf32> - iree.store_output(%1 : tensor<1xf32>, %arg1 : memref<1xf32>) - iree.return - } - } - } - } - iree.multi_arch_executable @main_ex_dispatch_8[8]() { - iree.executable[8](Unspecified) { - module { - func @main_entry_dispatch_8(%arg0: memref<1xf32>, %arg1: memref<1x10xf32>, %arg2: memref<1x10xf32>) - attributes {iree.executable.export, iree.executable.workload = dense<[10, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} { - %0 = iree.load_input(%arg0 : memref<1xf32>) : tensor<1xf32> - %1 = iree.load_input(%arg1 : memref<1x10xf32>) : tensor<1x10xf32> - %2 = "mhlo.broadcast_in_dim"(%0) {broadcast_dimensions = dense<0> : tensor<1xi64>, name = "broadcast.34"} : (tensor<1xf32>) -> tensor<1x10xf32> - %3 = divf %1, %2 : tensor<1x10xf32> - iree.store_output(%3 : tensor<1x10xf32>, %arg2 : memref<1x10xf32>) - iree.return - } - } - } - } - func @main(%arg0: memref<1x28x28x1xf32>) -> memref<1x10xf32> - attributes {iree.module.export} { - %0 = "iree_ll_seq.constant"() {value = dense<5.000000e-01> : tensor<784x128xf32>} : () -> memref<784x128xf32> - %1 = "iree_ll_seq.constant"() {value = dense<5.000000e-01> : tensor<128x10xf32>} : () -> memref<128x10xf32> - %2 = "iree_ll_seq.alloc_heap"() : () -> memref<1x784xf32> - iree_ll_seq.static_dispatch main_ex_dispatch_0::main_entry_dispatch_0[dense<[784, 1, 1]> : tensor<3xi32>](%arg0, %2) : (memref<1x28x28x1xf32>, memref<1x784xf32>) -> () - %3 = "iree_ll_seq.alloc_heap"() : () -> memref<1x128xf32> - iree_ll_seq.static_dispatch main_ex_dispatch_1::main_entry_dispatch_1[dense<[128, 1, 1]> : tensor<3xi32>](%2, %0, %3) : (memref<1x784xf32>, memref<784x128xf32>, memref<1x128xf32>) -> () - %4 = "iree_ll_seq.alloc_heap"() : () -> memref<1x128xf32> - iree_ll_seq.static_dispatch main_ex_dispatch_2::main_entry_dispatch_2[dense<[128, 1, 1]> : tensor<3xi32>](%3, %4) : (memref<1x128xf32>, memref<1x128xf32>) -> () - %5 = "iree_ll_seq.alloc_heap"() : () -> memref<1x10xf32> - iree_ll_seq.static_dispatch main_ex_dispatch_3::main_entry_dispatch_3[dense<[10, 1, 1]> : tensor<3xi32>](%4, %1, %5) : (memref<1x128xf32>, memref<128x10xf32>, memref<1x10xf32>) -> () - %6 = "iree_ll_seq.alloc_heap"() : () -> memref<1x10xf32> - iree_ll_seq.static_dispatch main_ex_dispatch_4::main_entry_dispatch_4[dense<[10, 1, 1]> : tensor<3xi32>](%5, %6) : (memref<1x10xf32>, memref<1x10xf32>) -> () - %7 = "iree_ll_seq.alloc_heap"() : () -> memref<1xf32> - iree_ll_seq.static_dispatch main_ex_dispatch_5::main_entry_dispatch_5[dense<1> : tensor<3xi32>](%6, %7) : (memref<1x10xf32>, memref<1xf32>) -> () - %8 = "iree_ll_seq.alloc_heap"() : () -> memref<1x10xf32> - iree_ll_seq.static_dispatch main_ex_dispatch_6::main_entry_dispatch_6[dense<[10, 1, 1]> : tensor<3xi32>](%6, %7, %8) : (memref<1x10xf32>, memref<1xf32>, memref<1x10xf32>) -> () - %9 = "iree_ll_seq.alloc_heap"() : () -> memref<1xf32> - iree_ll_seq.static_dispatch main_ex_dispatch_7::main_entry_dispatch_7[dense<1> : tensor<3xi32>](%8, %9) : (memref<1x10xf32>, memref<1xf32>) -> () - %10 = "iree_ll_seq.alloc_heap"() : () -> memref<1x10xf32> - iree_ll_seq.static_dispatch main_ex_dispatch_8::main_entry_dispatch_8[dense<[10, 1, 1]> : tensor<3xi32>](%9, %8, %10) : (memref<1xf32>, memref<1x10xf32>, memref<1x10xf32>) -> () - iree_ll_seq.return %10 : memref<1x10xf32> - } -} -``` - -**NOTE**: this is effectively compiling in -O0, which is why the buffers are not -aliased and some dispatch region fusing is not performed. As we get things going -we'll be adding simple optimizations that can operate on this IR to elide almost -all copies and externalize allocations to transient pooled memory. - -## Final IREE Module with SPIR-V - -TODO(benvanik): once reductions are done.
diff --git a/docs/using_colab.md b/docs/using_iree/using_colab.md similarity index 100% rename from docs/using_colab.md rename to docs/using_iree/using_colab.md
diff --git a/experimental/ModelBuilder/ModelRunner.h b/experimental/ModelBuilder/ModelRunner.h index 3539dbe..8afb37a 100644 --- a/experimental/ModelBuilder/ModelRunner.h +++ b/experimental/ModelBuilder/ModelRunner.h
@@ -116,7 +116,10 @@ } // Direct invocation based on MemRefType which automatically packs the data. template <typename... Args> - llvm::Error invoke(StringRef funcName, Args &... args) { + // TODO(suderman): Re-enable clang-format when new version migrates. + // clang-format off + llvm::Error invoke(StringRef funcName, Args &...args) { + // clang-format on const std::string adapterName = std::string("_mlir_ciface_") + funcName.str(); void *argsArray[] = {getData(args)...};
diff --git a/integrations/tensorflow/bindings/python/pyiree/tf/compiler/saved_model_test.py b/integrations/tensorflow/bindings/python/pyiree/tf/compiler/saved_model_test.py index 92f9c43..a7aba66 100644 --- a/integrations/tensorflow/bindings/python/pyiree/tf/compiler/saved_model_test.py +++ b/integrations/tensorflow/bindings/python/pyiree/tf/compiler/saved_model_test.py
@@ -12,10 +12,6 @@ # See the License for the specific language governing permissions and # limitations under the License. -from __future__ import absolute_import -from __future__ import division -from __future__ import print_function - import importlib import os import sys
diff --git a/integrations/tensorflow/bindings/python/pyiree/tf/compiler/signature_def_saved_model_test.py b/integrations/tensorflow/bindings/python/pyiree/tf/compiler/signature_def_saved_model_test.py index da94f6f..8a2e1cb 100644 --- a/integrations/tensorflow/bindings/python/pyiree/tf/compiler/signature_def_saved_model_test.py +++ b/integrations/tensorflow/bindings/python/pyiree/tf/compiler/signature_def_saved_model_test.py
@@ -12,10 +12,6 @@ # See the License for the specific language governing permissions and # limitations under the License. -from __future__ import absolute_import -from __future__ import division -from __future__ import print_function - import importlib import os import sys
diff --git a/integrations/tensorflow/bindings/python/pyiree/tf/support/tf_test_utils.py b/integrations/tensorflow/bindings/python/pyiree/tf/support/tf_test_utils.py index f032772..f2c6151 100644 --- a/integrations/tensorflow/bindings/python/pyiree/tf/support/tf_test_utils.py +++ b/integrations/tensorflow/bindings/python/pyiree/tf/support/tf_test_utils.py
@@ -264,24 +264,22 @@ return self.multi() reinitialized_modules = [ - tf_utils.CompiledModule.from_existing(module) - for module in compiled_backends.values() + module.create_reinitialized() for module in compiled_backends.values() ] return VirtualBackendsClass(*reinitialized_modules) -def compile_module(module_ctor, exported_names=()): - """SavedModelTestCase decorator that compiles a tf.Module. +def compile_module(module_class, exported_names=()): + """CompiledModuleTestCase decorator that compiles a tf.Module. A CompiledModule is created for each backend in --target_backends. They can be accessed individually via self.compiled_modules.backend_name or as a union via self.get_module(). Args: - module_ctor: tf.Module subclass or function which returns a tf.Module - subclass instance. + module_class: the tf.Module subclass to compile. exported_names: optional iterable of strings representing which of - module_ctor's functions to compile. If exported_names is empty all + module_class's functions to compile. If exported_names is empty all functions will be compiled. Returns: @@ -290,11 +288,11 @@ def decorator(cls): """Decorator Function.""" - if not issubclass(cls, SavedModelTestCase): + if not issubclass(cls, CompiledModuleTestCase): logging.exception( "The 'compile_module' decorator must be applied to a " - "SavedModelTestCase derived class, which %s is not.", cls) - cls._module_ctor = module_ctor + "CompiledModuleTestCase derived class, which %s is not.", cls) + cls._module_class = module_class cls._exported_names = exported_names return cls @@ -336,11 +334,11 @@ return backends -class SavedModelTestCase(tf.test.TestCase): - """Tests against a SavedModel.""" +class CompiledModuleTestCase(tf.test.TestCase): + """Compiles a tf.Module to multiple backends to test their correctness.""" # Will be initialized by the @compile_module decorator. - _module_ctor = None + _module_class = None _exported_names = () # Will be initialized in setUpClass to a dict of @@ -350,27 +348,33 @@ @classmethod def setUpClass(cls): super().setUpClass() - if cls._module_ctor is not None: - # Setup the debug directory for this test. Creates a global variable - # `global_debug_dir`. - _setup_test_debug_dir(test_name=cls.__name__) + if cls._module_class is None: + raise AttributeError( + "setUpClass was called but no module was specified. Specify a module " + "to compile via the @tf_test_utils.compile_module decorator.") - # Setup crash reproducer for the test. - crash_reproducer_path = os.path.join(global_debug_dir, "reproducer.mlir") - compiler.Context.default_crash_reproducer_path = crash_reproducer_path + # Setup the debug directory for this test. Creates a global variable + # `global_debug_dir`. + _setup_test_debug_dir(test_name=cls.__name__) - # Create a CompiledModule for each backend. - try: - backends = get_backends() - cls._compiled_backends_dict = {} - for backend in backends: - compiled_backend = tf_utils.CompiledModule.compile( - cls._module_ctor, backend, cls._exported_names, global_debug_dir) - cls._compiled_backends_dict[backend.name] = compiled_backend - finally: - # Disable crash reproducer (to avoid inadvertently overwriting this - # path on a subsequent interaction). - compiler.Context.default_crash_reproducer_path = None + # Setup crash reproducer for the test. + crash_reproducer_path = os.path.join(global_debug_dir, "reproducer.mlir") + compiler.Context.default_crash_reproducer_path = crash_reproducer_path + + # Create a CompiledModule for each backend. + try: + backends = get_backends() + cls._compiled_backends_dict = {} + for backend_info in backends: + compiled_backend = backend_info.CompiledModule(cls._module_class, + backend_info, + cls._exported_names, + global_debug_dir) + cls._compiled_backends_dict[backend_info.name] = compiled_backend + finally: + # Disable crash reproducer (to avoid inadvertently overwriting this + # path on a subsequent interaction). + compiler.Context.default_crash_reproducer_path = None @classmethod def tearDownClass(cls):
diff --git a/integrations/tensorflow/bindings/python/pyiree/tf/support/tf_utils.py b/integrations/tensorflow/bindings/python/pyiree/tf/support/tf_utils.py index 2752629..46a3785 100644 --- a/integrations/tensorflow/bindings/python/pyiree/tf/support/tf_utils.py +++ b/integrations/tensorflow/bindings/python/pyiree/tf/support/tf_utils.py
@@ -39,22 +39,32 @@ np.random.seed(seed) +def backends_to_str(target_backends): + """Creates a flattened and normalized string representing target_backends.""" + normalized_backends = [] + for backend in target_backends: + # Remove unusual characters and ensure names don't end or start in "_". + backend = re.sub("[^0-9a-zA-Z_]+", "_", backend) + normalized_backends.append(backend.strip("_")) + return "__".join(normalized_backends) + + def compile_tf_module(tf_module, target_backends=(), exported_names=(), artifacts_dir=None): """Compiles a TensorFlow tf.Module and optionally saves compilation artifacts. - The artifact this creates is not callable. See IreeCompiledModule.compile(...) - for an API that returns a module that can be called without any further steps. + The artifact this creates is not callable. See IreeCompiledModule for an API + that returns a module that can be called without any further steps. If artifacts_dir is provided then the following artifacts will be saved: saved_model: A TF SavedModel directory containing the files used translate the tf.Module into an IREE module. - tf_input__backends.mlir: + tf_input.mlir: MLIR for the module in TF's input dialect. - iree_input__backends.mlir: + iree_input.mlir: The MLIR above translated to IREE via compiler.TF_IMPORT_PASS_PIPELINE. compiled__backends.vmfb: A VM FlatBuffer compiled to the target backends from the IREE MLIR above. @@ -77,14 +87,6 @@ # We break up the compilation here so we can save intermediary artifacts. compiler_context = compiler.Context() - if artifacts_dir is not None: - normalized_backends = [] - for backend in target_backends: - # Remove unusual characters and ensure names don't end or start in "_". - backend = re.sub("[^0-9a-zA-Z_]+", "_", backend) - normalized_backends.append(backend.strip("_")) - backends_string = "__".join(normalized_backends) - # Convert the tf_module into raw TF input MLIR. compiler_module = compiler.tf_load_saved_model( sm_path, @@ -93,8 +95,7 @@ pass_pipeline=()) if artifacts_dir is not None: - tf_mlir_path = os.path.join(artifacts_dir, - f"tf_input__{backends_string}.mlir") + tf_mlir_path = os.path.join(artifacts_dir, "tf_input.mlir") logging.info("Saving raw TF input MLIR to: %s", tf_mlir_path) with open(tf_mlir_path, "w") as f: f.write(compiler_module.to_asm()) @@ -103,16 +104,15 @@ compiler_module.run_pass_pipeline(compiler.TF_IMPORT_PASS_PIPELINE) if artifacts_dir is not None: - iree_mlir_path = os.path.join(artifacts_dir, - f"iree_input__{backends_string}.mlir") + iree_mlir_path = os.path.join(artifacts_dir, "iree_input.mlir") logging.info("Saving IREE input MLIR to: %s", iree_mlir_path) with open(iree_mlir_path, "w") as f: f.write(compiler_module.to_asm()) compiled_module = compiler_module.compile(target_backends=target_backends) if artifacts_dir is not None: - compiled_path = os.path.join(artifacts_dir, - f"compiled__{backends_string}.vmfb") + compiled_name = f"compiled__{backends_to_str(target_backends)}.vmfb" + compiled_path = os.path.join(artifacts_dir, compiled_name) logging.info("Saving compiled IREE module to: %s", compiled_path) with open(compiled_path, "wb") as f: f.write(compiled_module) @@ -133,97 +133,67 @@ class CompiledModule(object): - """Base class for the TF and IREE compiled module facades.""" + """Base class for the TF and IREE compiled modules.""" - @staticmethod - def compile(constructor, backend_info, exported_names=(), artifacts_dir=None): - """Compile a tf.Module using the CompiledModule subclass in backend_info. - - Args: - constructor: a tf.Module subclass or function which returns a tf.Module - subclass instance. - backend_info: an element of BackendInfo corresponding to the backend to - compile to. If a TF 'backend' is provided then the module is wrapped in - a TfCompiledModule. - exported_names: an optional iterable of strings representing which of the - tf.Module's functions to compile. If exported_names is empty all - functions will be compiled. - artifacts_dir: an optional path to save compilation artifacts to. - """ - compile = backend_info.CompiledModule.compile - return compile(constructor, backend_info, exported_names, artifacts_dir) - - @staticmethod - def from_existing(module): - """Duplicates 'module' with the tf.Module's state without recompiling.""" - # Use the backend_info attr to determine which subclass' constructor to use. - from_existing = module._backend_info.CompiledModule.from_existing - return from_existing(module) - - def __init__(self, constructor, backend_info, exported_names, artifacts_dir): - """Default constructor – use `compile` or `from_existing` instead.""" - self._constructor = constructor + def __init__(self, module_class, backend_info, exported_names, artifacts_dir): + """Shared base constructor – not useful on its own.""" + self._module_class = module_class self._backend_info = backend_info self._exported_names = exported_names self._artifacts_dir = artifacts_dir + def create_reinitialized(self): + """Duplicates this module with its initial state without recompiling.""" + raise NotImplementedError() + class IreeCompiledModule(CompiledModule): """Iree compiled module.""" - @staticmethod - def compile(constructor, backend_info, exported_names=(), artifacts_dir=None): + def __init__(self, + module_class, + backend_info, + exported_names=[], + artifacts_dir=None, + _create_reinitialized_args=None): """Compile a tf.Module to the target backend in backend_info. Args: - constructor: a tf.Module subclass or function which returns a tf.Module - subclass instance. + module_class: the tf.Module subclass to compile. backend_info: an element of BackendInfo corresponding to the IREE backend to compile to. exported_names: an optional iterable of strings representing which of the - tf.Module's functions to compile. If exported_names is empty all + module_class's functions to compile. If exported_names is empty all functions will be compiled. artifacts_dir: an optional path to save compilation artifacts to. """ - return IreeCompiledModule(constructor, backend_info, exported_names, - artifacts_dir) + super().__init__(module_class, backend_info, exported_names, artifacts_dir) - @staticmethod - def from_existing(module): - """Duplicates 'module' with the tf.Module's state without recompiling.""" - default_args = [ - module._constructor, module._backend_info, module._exported_names, - module._artifacts_dir - ] - from_existing_args = [module._module_blob, module._module, module._config] - return IreeCompiledModule(*default_args, from_existing_args) - - def __init__(self, - constructor, - backend_info, - exported_names, - artifacts_dir, - _from_existing_args=None): - """Default constructor – use `compile` or `from_existing` instead.""" - super().__init__(constructor, backend_info, exported_names, artifacts_dir) - - if _from_existing_args is None: - # Called from IreeCompiledModule.compile(...) + if _create_reinitialized_args is None: self._module_blob = compile_tf_module( - tf_module=constructor(), + tf_module=module_class(), target_backends=backend_info.iree_compiler_targets, exported_names=exported_names, artifacts_dir=artifacts_dir) self._module = rt.VmModule.from_flatbuffer(self._module_blob) self._config = rt.Config(driver_name=backend_info.iree_driver) else: - # Called from IreeCompiledModule.from_existing(module) - self._module_blob, self._module, self._config = _from_existing_args + # Called from self.create_reinitialized() + self._module_blob, self._module, self._config = _create_reinitialized_args # Holds all of the module's mutable state. self._context = rt.SystemContext( modules=[self._module], config=self._config) + def create_reinitialized(self): + """Duplicates this module with its initial state without recompiling.""" + default_args = [ + self._module_class, self._backend_info, self._exported_names, + self._artifacts_dir + ] + create_reinitialized_args = [self._module_blob, self._module, self._config] + return IreeCompiledModule(*default_args, create_reinitialized_args) + def __getattr__(self, attr): # Try to resolve it as a function. m = self._context.modules[self._module.name] @@ -249,36 +219,29 @@ normalize TensorFlow's output to Numpy. """ - @staticmethod - def compile(constructor, backend_info, exported_names=(), artifacts_dir=None): + def __init__(self, + module_class, + backend_info, + exported_names=[], + artifacts_dir=None): """Wrap a tf.Module in a TFCompiledModule facade. Args: - constructor: a tf.Module subclass or function which returns a tf.Module - subclass instance. + module_class: the tf.Module subclass to 'compile'. backend_info: one of the 'tf*' elements in BackendInfo. - exported_names: an optional iterable of strings representing the which of - the tf.Module's functions should be callable. If exported_names is empty - then all functions are callable. + exported_names: an optional iterable of strings representing which of the + module_class's functions should be callable. If exported_names is empty + then all functions will be callable. artifacts_dir: an optional path to save compilation artifacts to. Has no effect for this subclass as nothing is compiled. """ - return TfCompiledModule(constructor, backend_info, exported_names, - artifacts_dir) + super().__init__(module_class, backend_info, exported_names, artifacts_dir) + self._tf_module = module_class() - @staticmethod - def from_existing(module): - """Duplicates 'module's facade with the starting state of constructor.""" - duplicate_module = TfCompiledModule(module._constructor, - module._backend_info, - module._exported_names, - module._artifacts_dir) - return duplicate_module - - def __init__(self, constructor, backend_info, exported_names, artifacts_dir): - """Default constructor – use `compile` or `from_existing` instead.""" - super().__init__(constructor, backend_info, exported_names, artifacts_dir) - self._tf_module = constructor() + def create_reinitialized(self): + """Duplicates this module with the starting state of module_class.""" + return TfCompiledModule(self._module_class, self._backend_info, + self._exported_names, self._artifacts_dir) def __getattr__(self, attr): # Try to resolve it as a function.
diff --git a/integrations/tensorflow/bindings/python/pyiree/tf/support/tf_utils_test.py b/integrations/tensorflow/bindings/python/pyiree/tf/support/tf_utils_test.py index dcc3aec..fe25f97 100644 --- a/integrations/tensorflow/bindings/python/pyiree/tf/support/tf_utils_test.py +++ b/integrations/tensorflow/bindings/python/pyiree/tf/support/tf_utils_test.py
@@ -17,6 +17,7 @@ import os import tempfile +from absl import logging from absl.testing import parameterized from pyiree.tf.support import tf_utils import tensorflow as tf @@ -52,7 +53,7 @@ }, { 'testcase_name': 'multiple_backends', - 'target_backends': ['vmla', 'llvm'], + 'target_backends': ['vmla', 'llvm-ir'], }, ]) def test_artifact_saving(self, target_backends): @@ -64,13 +65,13 @@ artifacts_dir=artifacts_dir) artifacts_to_check = [ - 'saved_model', - f'tf_input__{"__".join(target_backends)}.mlir', - f'iree_input__{"__".join(target_backends)}.mlir', - f'compiled__{"__".join(target_backends)}.vmfb', + 'saved_model', 'tf_input.mlir', 'iree_input.mlir', + f'compiled__{tf_utils.backends_to_str(target_backends)}.vmfb', ] for artifact in artifacts_to_check: - self.assertTrue(os.path.exists(os.path.join(artifacts_dir, artifact))) + artifact_path = os.path.join(artifacts_dir, artifact) + logging.info('Checking path: %s', artifact_path) + self.assertTrue(os.path.exists(artifact_path)) @parameterized.named_parameters([ { @@ -83,15 +84,15 @@ }, ]) def test_unaltered_state(self, backend_name): - info = tf_utils.BackendInfo.ALL[backend_name] - module = tf_utils.CompiledModule.compile(StatefulCountingModule, info) + backend_info = tf_utils.BackendInfo.ALL[backend_name] + module = backend_info.CompiledModule(StatefulCountingModule, backend_info) # Test that incrementing works properly. self.assertEqual([0.], module.get_count()) module.increment() self.assertEqual([1.], module.get_count()) - reinitialized_module = tf_utils.CompiledModule.from_existing(module) + reinitialized_module = module.create_reinitialized() # Test reinitialization. self.assertEqual([0.], reinitialized_module.get_count()) # Test independent state.
diff --git a/integrations/tensorflow/bindings/python/pyiree/xla/compiler/xla_module_proto_test.py b/integrations/tensorflow/bindings/python/pyiree/xla/compiler/xla_module_proto_test.py index eac1288..02aeb39 100644 --- a/integrations/tensorflow/bindings/python/pyiree/xla/compiler/xla_module_proto_test.py +++ b/integrations/tensorflow/bindings/python/pyiree/xla/compiler/xla_module_proto_test.py
@@ -12,10 +12,6 @@ # See the License for the specific language governing permissions and # limitations under the License. -from __future__ import absolute_import -from __future__ import division -from __future__ import print_function - from absl.testing import absltest import numpy as np from pyiree.xla import compiler
diff --git a/integrations/tensorflow/e2e/README.md b/integrations/tensorflow/e2e/README.md index 341aa56..04604d8 100644 --- a/integrations/tensorflow/e2e/README.md +++ b/integrations/tensorflow/e2e/README.md
@@ -26,10 +26,10 @@ ## Compiling `tf.Module`s Compatible TensorFlow modules can be compiled to specific IREE backends using -`IreeCompiledModule.compile(...)`. This also optionally saves compilation -artifacts to a specified directory. These artifacts include: MLIR across various -lowerings, a TensorFlow SavedModel, and the compiled VM FlatBuffer. A basic -example of creating and calling an `IreeCompiledModule` can be found in +`IreeCompiledModule`. This also optionally saves compilation artifacts to a +specified directory. These artifacts include: MLIR across various lowerings, a +TensorFlow SavedModel, and the compiled VM FlatBuffer. A basic example of +creating and calling an `IreeCompiledModule` can be found in [`tf_utils_test.py`](https://github.com/google/iree/blob/main/integrations/tensorflow/bindings/python/pyiree/tf/support/tf_utils_test.py) When using Keras models or tf.Modules with functions that IREE can't compile, @@ -38,7 +38,7 @@ ```python from pyiree.tf.support import tf_utils vmla_module = tf_utils.IreeCompiledModule( - constructor=KerasTFModuleClass, + module_class=KerasTFModuleClass, backend_info=tf_utils.BackendInfo.ALL['iree_vmla'], exported_names=['predict']) vmla_module.predict(...)
diff --git a/integrations/tensorflow/e2e/batch_norm_test.py b/integrations/tensorflow/e2e/batch_norm_test.py index f9f8d8c..75de16d 100644 --- a/integrations/tensorflow/e2e/batch_norm_test.py +++ b/integrations/tensorflow/e2e/batch_norm_test.py
@@ -39,7 +39,7 @@ @tf_test_utils.compile_module(BatchNormModule) -class BatchNormTest(tf_test_utils.SavedModelTestCase): +class BatchNormTest(tf_test_utils.CompiledModuleTestCase): def test_batch_norm_inference(self): np.random.seed(12345)
diff --git a/integrations/tensorflow/e2e/broadcasting_test.py b/integrations/tensorflow/e2e/broadcasting_test.py index cde2fd6..74880bd 100644 --- a/integrations/tensorflow/e2e/broadcasting_test.py +++ b/integrations/tensorflow/e2e/broadcasting_test.py
@@ -29,7 +29,7 @@ @tf_test_utils.compile_module(BroadcastingModule) -class BroadcastingTest(tf_test_utils.SavedModelTestCase): +class BroadcastingTest(tf_test_utils.CompiledModuleTestCase): def test_add_same_shape(self): m = self.get_module()
diff --git a/integrations/tensorflow/e2e/concat_test.py b/integrations/tensorflow/e2e/concat_test.py index a9f9759..b7a348c 100644 --- a/integrations/tensorflow/e2e/concat_test.py +++ b/integrations/tensorflow/e2e/concat_test.py
@@ -51,7 +51,7 @@ @tf_test_utils.compile_module(ConcatOpsModule) -class ConcatOpsTest(tf_test_utils.SavedModelTestCase): +class ConcatOpsTest(tf_test_utils.CompiledModuleTestCase): def test_concat_zero_dim(self): tf_utils.set_random_seed()
diff --git a/integrations/tensorflow/e2e/control_flow_test.py b/integrations/tensorflow/e2e/control_flow_test.py index 0223e8c..0c25fd6 100644 --- a/integrations/tensorflow/e2e/control_flow_test.py +++ b/integrations/tensorflow/e2e/control_flow_test.py
@@ -12,10 +12,6 @@ # See the License for the specific language governing permissions and # limitations under the License. -from __future__ import absolute_import -from __future__ import division -from __future__ import print_function - import numpy from pyiree.tf.support import tf_test_utils import tensorflow.compat.v2 as tf @@ -39,7 +35,7 @@ @tf_test_utils.compile_module(ControlFlowModule) -class ControlFlowTest(tf_test_utils.SavedModelTestCase): +class ControlFlowTest(tf_test_utils.CompiledModuleTestCase): def test_short_sequence(self): input_array = numpy.array(9., dtype=numpy.float32)
diff --git a/integrations/tensorflow/e2e/conv_test.py b/integrations/tensorflow/e2e/conv_test.py index f72b11d..61c46cf 100644 --- a/integrations/tensorflow/e2e/conv_test.py +++ b/integrations/tensorflow/e2e/conv_test.py
@@ -99,7 +99,7 @@ @tf_test_utils.compile_module(Conv2dModule) -class ConvTest(tf_test_utils.SavedModelTestCase): +class ConvTest(tf_test_utils.CompiledModuleTestCase): def test_id_batch_size_1(self): i = np.arange(20, dtype=np.float32).reshape([1, 4, 5, 1])
diff --git a/integrations/tensorflow/e2e/depth_conv_test.py b/integrations/tensorflow/e2e/depth_conv_test.py index cdf4d1e..1e8a002 100644 --- a/integrations/tensorflow/e2e/depth_conv_test.py +++ b/integrations/tensorflow/e2e/depth_conv_test.py
@@ -39,7 +39,7 @@ @tf_test_utils.compile_module(Conv2dModule) -class ConvTest(tf_test_utils.SavedModelTestCase): +class ConvTest(tf_test_utils.CompiledModuleTestCase): def test_batched_feature_unpadded(self): i = np.arange(80, dtype=np.float32).reshape([2, 4, 5, 2])
diff --git a/integrations/tensorflow/e2e/dynamic_mlp_relu_test.py b/integrations/tensorflow/e2e/dynamic_mlp_relu_test.py index 04de603..64c51e9 100644 --- a/integrations/tensorflow/e2e/dynamic_mlp_relu_test.py +++ b/integrations/tensorflow/e2e/dynamic_mlp_relu_test.py
@@ -66,7 +66,7 @@ @tf_test_utils.compile_module(Mlp, exported_names=["predict"]) -class DynamicMlpTest(tf_test_utils.SavedModelTestCase): +class DynamicMlpTest(tf_test_utils.CompiledModuleTestCase): def test_dynamic_batch(self): m = self.get_module()
diff --git a/integrations/tensorflow/e2e/dynamic_mlp_test.py b/integrations/tensorflow/e2e/dynamic_mlp_test.py index 66f3c06..72d7f1f 100644 --- a/integrations/tensorflow/e2e/dynamic_mlp_test.py +++ b/integrations/tensorflow/e2e/dynamic_mlp_test.py
@@ -62,7 +62,7 @@ @tf_test_utils.compile_module(Mlp, exported_names=["predict"]) -class DynamicMlpTest(tf_test_utils.SavedModelTestCase): +class DynamicMlpTest(tf_test_utils.CompiledModuleTestCase): def test_dynamic_batch(self): m = self.get_module()
diff --git a/integrations/tensorflow/e2e/explicit_backend_test.py b/integrations/tensorflow/e2e/explicit_backend_test.py index 903b34c..bdcdd79 100644 --- a/integrations/tensorflow/e2e/explicit_backend_test.py +++ b/integrations/tensorflow/e2e/explicit_backend_test.py
@@ -30,7 +30,7 @@ @tf_test_utils.compile_module(SimpleArithmeticModule) -class ExplicitBackendTest(tf_test_utils.SavedModelTestCase): +class ExplicitBackendTest(tf_test_utils.CompiledModuleTestCase): def test_explicit(self): a = np.array([1., 2., 3., 4.], dtype=np.float32)
diff --git a/integrations/tensorflow/e2e/fill_test.py b/integrations/tensorflow/e2e/fill_test.py index 8ef96a9..8d912a4 100644 --- a/integrations/tensorflow/e2e/fill_test.py +++ b/integrations/tensorflow/e2e/fill_test.py
@@ -31,7 +31,7 @@ @tf_test_utils.compile_module(FillModule) -class FillTest(tf_test_utils.SavedModelTestCase): +class FillTest(tf_test_utils.CompiledModuleTestCase): def test_fill(self): dims = np.array([2, 3], dtype=np.int32)
diff --git a/integrations/tensorflow/e2e/gather_test.py b/integrations/tensorflow/e2e/gather_test.py index 8d2a0ba..67f5acf 100644 --- a/integrations/tensorflow/e2e/gather_test.py +++ b/integrations/tensorflow/e2e/gather_test.py
@@ -49,7 +49,7 @@ @tf_test_utils.compile_module(GatherModule) -class GatherTest(tf_test_utils.SavedModelTestCase): +class GatherTest(tf_test_utils.CompiledModuleTestCase): def test_gather_axis0_scalar(self): indices = np.array(2, dtype=np.int32)
diff --git a/integrations/tensorflow/e2e/keras/lstm_static_test.py b/integrations/tensorflow/e2e/keras/lstm_static_test.py index 0d34d97..fb7a58c 100644 --- a/integrations/tensorflow/e2e/keras/lstm_static_test.py +++ b/integrations/tensorflow/e2e/keras/lstm_static_test.py
@@ -27,21 +27,23 @@ INPUT_SHAPE = [NUM_BATCH, NUM_TIMESTEPS, NUM_UNITS] -def lstm_module(): - tf_utils.set_random_seed() - inputs = tf.keras.layers.Input(batch_size=NUM_BATCH, shape=INPUT_SHAPE[1:]) - outputs = tf.keras.layers.LSTM(units=NUM_UNITS, return_sequences=True)(inputs) - model = tf.keras.Model(inputs, outputs) - module = tf.Module() - module.m = model - module.predict = tf.function( - input_signature=[tf.TensorSpec(INPUT_SHAPE, tf.float32)])( - model.call) - return module +class LstmStatic(tf.Module): + + def __init__(self): + super(LstmStatic, self).__init__() + tf_utils.set_random_seed() + inputs = tf.keras.layers.Input(batch_size=NUM_BATCH, shape=INPUT_SHAPE[1:]) + outputs = tf.keras.layers.LSTM( + units=NUM_UNITS, return_sequences=True)( + inputs) + self.m = tf.keras.Model(inputs, outputs) + self.predict = tf.function( + input_signature=[tf.TensorSpec(INPUT_SHAPE, tf.float32)])( + self.m.call) -@tf_test_utils.compile_module(lstm_module, exported_names=["predict"]) -class LstmTest(tf_test_utils.SavedModelTestCase): +@tf_test_utils.compile_module(LstmStatic, exported_names=["predict"]) +class LstmTest(tf_test_utils.CompiledModuleTestCase): def test_lstm(self): m = self.get_module()
diff --git a/integrations/tensorflow/e2e/keras/lstm_test.py b/integrations/tensorflow/e2e/keras/lstm_test.py index 671c31b..9409d04 100644 --- a/integrations/tensorflow/e2e/keras/lstm_test.py +++ b/integrations/tensorflow/e2e/keras/lstm_test.py
@@ -24,21 +24,23 @@ INPUT_SHAPE = [None, None, NUM_UNITS] -def lstm_module(): - tf_utils.set_random_seed() - inputs = tf.keras.layers.Input(batch_size=None, shape=INPUT_SHAPE[1:]) - outputs = tf.keras.layers.LSTM(units=NUM_UNITS, return_sequences=True)(inputs) - model = tf.keras.Model(inputs, outputs) - module = tf.Module() - module.m = model - module.predict = tf.function( - input_signature=[tf.TensorSpec(INPUT_SHAPE, tf.float32)])( - model.call) - return module +class Lstm(tf.Module): + + def __init__(self): + super(Lstm, self).__init__() + tf_utils.set_random_seed() + inputs = tf.keras.layers.Input(batch_size=None, shape=INPUT_SHAPE[1:]) + outputs = tf.keras.layers.LSTM( + units=NUM_UNITS, return_sequences=True)( + inputs) + self.m = tf.keras.Model(inputs, outputs) + self.predict = tf.function( + input_signature=[tf.TensorSpec(INPUT_SHAPE, tf.float32)])( + self.m.call) -@tf_test_utils.compile_module(lstm_module, exported_names=["predict"]) -class LstmTest(tf_test_utils.SavedModelTestCase): +@tf_test_utils.compile_module(Lstm, exported_names=["predict"]) +class LstmTest(tf_test_utils.CompiledModuleTestCase): def test_lstm(self): m = self.get_module()
diff --git a/integrations/tensorflow/e2e/keras/train/model_train_test.py b/integrations/tensorflow/e2e/keras/train/model_train_test.py index 6675956..e30bd57 100644 --- a/integrations/tensorflow/e2e/keras/train/model_train_test.py +++ b/integrations/tensorflow/e2e/keras/train/model_train_test.py
@@ -77,7 +77,7 @@ @tf_test_utils.compile_module( ModelTrain.CreateModule, exported_names=["TrainStep"]) -class ModelTrainTest(tf_test_utils.SavedModelTestCase): +class ModelTrainTest(tf_test_utils.CompiledModuleTestCase): def generate_regression_data(self, size=8): x = np.arange(size) - size // 2
diff --git a/integrations/tensorflow/e2e/keras/train_vision_models_on_cifar.py b/integrations/tensorflow/e2e/keras/train_vision_models_on_cifar.py index 0a79452..887bfb8 100644 --- a/integrations/tensorflow/e2e/keras/train_vision_models_on_cifar.py +++ b/integrations/tensorflow/e2e/keras/train_vision_models_on_cifar.py
@@ -13,9 +13,6 @@ # See the License for the specific language governing permissions and # limitations under the License. """Train vision models on CIFAR10.""" -from __future__ import absolute_import -from __future__ import division -from __future__ import print_function import os from absl import flags
diff --git a/integrations/tensorflow/e2e/keras/vision_model_test.py b/integrations/tensorflow/e2e/keras/vision_model_test.py index 1804739..54374ff 100644 --- a/integrations/tensorflow/e2e/keras/vision_model_test.py +++ b/integrations/tensorflow/e2e/keras/vision_model_test.py
@@ -14,6 +14,8 @@ # limitations under the License. """Test all applications models in Keras.""" import os + +from absl import app from absl import flags import numpy as np from pyiree.tf.support import tf_test_utils @@ -72,74 +74,84 @@ } -def get_input_shape(data, model): - if data == 'imagenet': - if (model == 'InceptionV3' or model == 'Xception' or - model == 'InceptionResNetV2'): +def get_input_shape(): + if FLAGS.data == 'imagenet': + if FLAGS.model in ['InceptionV3', 'Xception', 'InceptionResNetV2']: return (1, 299, 299, 3) - elif model == 'NASNetLarge': + elif FLAGS.model == 'NASNetLarge': return (1, 331, 331, 3) else: return (1, 224, 224, 3) - elif data == 'cifar10': + elif FLAGS.data == 'cifar10': return (1, 32, 32, 3) else: - raise ValueError('Not supported data ', data) + raise ValueError(f'Data not supported: {FLAGS.data}') -def models(): - tf.keras.backend.set_learning_phase(False) +def load_cifar10_weights(model): + file_name = 'cifar10' + FLAGS.model + # get_file will download the model weights from a publicly available folder, + # save them to cache_dir=~/.keras/models/ and return a path to them. + url = os.path.join( + FLAGS.url, f'cifar10_include_top_{FLAGS.include_top}_{FLAGS.model}.h5') + weights_path = tf.keras.utils.get_file(file_name, url) + model.load_weights(weights_path) + return model + + +def initialize_model(): tf_utils.set_random_seed() + tf.keras.backend.set_learning_phase(False) - input_shape = get_input_shape(FLAGS.data, FLAGS.model) - # keras model receives images size as input, - # where batch size is not specified - by default it is dynamic - if FLAGS.model in APP_MODELS: - weights = 'imagenet' if FLAGS.data == 'imagenet' else None + # Keras applications models receive input shapes without a batch dimension, as + # the batch size is dynamic by default. This selects just the image size. + input_shape = get_input_shape()[1:] - # if weights == 'imagenet' it will load weights from external tf.keras URL - model = APP_MODELS[FLAGS.model]( - weights=weights, - include_top=FLAGS.include_top, - input_shape=input_shape[1:]) + # If weights == 'imagenet', the model will load the appropriate weights from + # an external tf.keras URL. + weights = 'imagenet' if FLAGS.data == 'imagenet' else None - if FLAGS.data == 'cifar10' and FLAGS.url: - file_name = 'cifar10' + FLAGS.model - # it will download model weights from publically available folder: PATH - # and save it to cache_dir=~/.keras and return path to it - weights_path = tf.keras.utils.get_file( - file_name, - os.path.join( - FLAGS.url, - 'cifar10_include_top_{}_{}'.format(FLAGS.include_top, - FLAGS.model + '.h5'))) + model = APP_MODELS[FLAGS.model]( + weights=weights, include_top=FLAGS.include_top, input_shape=input_shape) - model.load_weights(weights_path) - else: - raise ValueError('Unsupported model', FLAGS.model) - - module = tf.Module() - module.m = model - # specify input size with static batch size - # TODO(b/142948097): with support of dynamic shape - # replace input_shape by model.input_shape, so batch size will be dynamic (-1) - module.predict = tf.function(input_signature=[tf.TensorSpec(input_shape)])( - model.call) - return module + if FLAGS.data == 'cifar10' and FLAGS.url: + model = load_cifar10_weights(model) + return model -@tf_test_utils.compile_module(models, exported_names=['predict']) -class AppTest(tf_test_utils.SavedModelTestCase): +class VisionModule(tf.Module): + + def __init__(self): + super(VisionModule, self).__init__() + self.m = initialize_model() + # Specify input shape with a static batch size. + # TODO(b/142948097): Add support for dynamic shapes in SPIR-V lowering. + # Replace input_shape with m.input_shape to make the batch size dynamic. + self.predict = tf.function( + input_signature=[tf.TensorSpec(get_input_shape())])( + self.m.call) + + +@tf_test_utils.compile_module(VisionModule, exported_names=['predict']) +class AppTest(tf_test_utils.CompiledModuleTestCase): def test_application(self): - input_shape = get_input_shape(FLAGS.data, FLAGS.model) - input_data = np.random.rand(np.prod(np.array(input_shape))).astype( - np.float32) - input_data = input_data.reshape(input_shape) + input_data = np.random.rand(*get_input_shape()).astype(np.float32) self.get_module().predict(input_data).print().assert_all_close(atol=1e-6) -if __name__ == '__main__': +def main(argv): + del argv # Unused if hasattr(tf, 'enable_v2_behavior'): tf.enable_v2_behavior() + + if FLAGS.model not in APP_MODELS: + raise ValueError(f'Unsupported model: {FLAGS.model}') + # Override VisionModule's __name__ to be more specific. + VisionModule.__name__ = FLAGS.model + tf.test.main() + + +if __name__ == '__main__': + app.run(main)
diff --git a/integrations/tensorflow/e2e/linspace_test.py b/integrations/tensorflow/e2e/linspace_test.py index d326db5..aa49f5b 100644 --- a/integrations/tensorflow/e2e/linspace_test.py +++ b/integrations/tensorflow/e2e/linspace_test.py
@@ -34,7 +34,7 @@ @tf_test_utils.compile_module(LinSpaceModule) -class LinspaceTest(tf_test_utils.SavedModelTestCase): +class LinspaceTest(tf_test_utils.CompiledModuleTestCase): def test_linspace(self): start = np.array(10., dtype=np.float32)
diff --git a/integrations/tensorflow/e2e/mandelbrot_test.py b/integrations/tensorflow/e2e/mandelbrot_test.py index 4886b7a..2b3a8d9 100644 --- a/integrations/tensorflow/e2e/mandelbrot_test.py +++ b/integrations/tensorflow/e2e/mandelbrot_test.py
@@ -12,10 +12,6 @@ # See the License for the specific language governing permissions and # limitations under the License. -from __future__ import absolute_import -from __future__ import division -from __future__ import print_function - from pyiree.tf.support import tf_test_utils import tensorflow.compat.v2 as tf @@ -95,7 +91,7 @@ @tf_test_utils.compile_module(MandelbrotModule) -class MandelbrotTest(tf_test_utils.SavedModelTestCase): +class MandelbrotTest(tf_test_utils.CompiledModuleTestCase): def test_mandelbrot(self): mandelbrot = self.get_module()
diff --git a/integrations/tensorflow/e2e/math_test.py b/integrations/tensorflow/e2e/math_test.py index b27d1d1..a33ac7c 100644 --- a/integrations/tensorflow/e2e/math_test.py +++ b/integrations/tensorflow/e2e/math_test.py
@@ -39,7 +39,7 @@ @tf_test_utils.compile_module(MathModule) -class MathTest(tf_test_utils.SavedModelTestCase): +class MathTest(tf_test_utils.CompiledModuleTestCase): def test_abs(self): a = np.array([-0.5, 0.0, 0.5, 1.0], dtype=np.float32)
diff --git a/integrations/tensorflow/e2e/matrix_ops_test.py b/integrations/tensorflow/e2e/matrix_ops_test.py index b29a198..d04ce3a 100644 --- a/integrations/tensorflow/e2e/matrix_ops_test.py +++ b/integrations/tensorflow/e2e/matrix_ops_test.py
@@ -71,7 +71,7 @@ @tf_test_utils.compile_module(MatrixOpsModule) -class MatrixOpsTest(tf_test_utils.SavedModelTestCase): +class MatrixOpsTest(tf_test_utils.CompiledModuleTestCase): def test_basic_matmul(self): m = self.get_module()
diff --git a/integrations/tensorflow/e2e/resource_ops_test.py b/integrations/tensorflow/e2e/resource_ops_test.py index 1d703c0..8daa6cf 100644 --- a/integrations/tensorflow/e2e/resource_ops_test.py +++ b/integrations/tensorflow/e2e/resource_ops_test.py
@@ -29,7 +29,7 @@ @tf_test_utils.compile_module(ResourcesOpsModule) -class ResourcesOpsTest(tf_test_utils.SavedModelTestCase): +class ResourcesOpsTest(tf_test_utils.CompiledModuleTestCase): def test_add_assign(self): result = self.get_module().add_assign(np.array(9., dtype=np.float32))
diff --git a/integrations/tensorflow/e2e/ring_buffer_test.py b/integrations/tensorflow/e2e/ring_buffer_test.py index ea48711..3af1502 100644 --- a/integrations/tensorflow/e2e/ring_buffer_test.py +++ b/integrations/tensorflow/e2e/ring_buffer_test.py
@@ -179,7 +179,7 @@ @tf_test_utils.compile_module( StatefulRingBufferModule, exported_names=["predict"]) -class StatefulRingBufferTest(tf_test_utils.SavedModelTestCase): +class StatefulRingBufferTest(tf_test_utils.CompiledModuleTestCase): def test_stateful_ringbuffer(self): input1 = np.array([[1.0, 2.0]], dtype=np.float32)
diff --git a/integrations/tensorflow/e2e/scatter_update_test.py b/integrations/tensorflow/e2e/scatter_update_test.py index cdd3277..ab5ab91 100644 --- a/integrations/tensorflow/e2e/scatter_update_test.py +++ b/integrations/tensorflow/e2e/scatter_update_test.py
@@ -49,7 +49,7 @@ @tf_test_utils.compile_module(ScatterUpdateModule) -class ScatterUpdateTest(tf_test_utils.SavedModelTestCase): +class ScatterUpdateTest(tf_test_utils.CompiledModuleTestCase): def test_scatter_update_1D(self): tensor = tf.ones([8], dtype=tf.int32)
diff --git a/integrations/tensorflow/e2e/simple_arithmetic_test.py b/integrations/tensorflow/e2e/simple_arithmetic_test.py index 0c5941d..d3ea327 100644 --- a/integrations/tensorflow/e2e/simple_arithmetic_test.py +++ b/integrations/tensorflow/e2e/simple_arithmetic_test.py
@@ -37,7 +37,7 @@ @tf_test_utils.compile_module(SimpleArithmeticModule) -class SimpleArithmeticTest(tf_test_utils.SavedModelTestCase): +class SimpleArithmeticTest(tf_test_utils.CompiledModuleTestCase): def test_simple_mul(self): a = np.array([1., 2., 3., 4.], dtype=np.float32)
diff --git a/integrations/tensorflow/e2e/simple_stateful_test.py b/integrations/tensorflow/e2e/simple_stateful_test.py index 45eba4f..24dd23e 100644 --- a/integrations/tensorflow/e2e/simple_stateful_test.py +++ b/integrations/tensorflow/e2e/simple_stateful_test.py
@@ -33,7 +33,7 @@ @tf_test_utils.compile_module(Stateful) -class StatefulTest(tf_test_utils.SavedModelTestCase): +class StatefulTest(tf_test_utils.CompiledModuleTestCase): def test_stateful(self): m = self.get_module()
diff --git a/integrations/tensorflow/e2e/sliding_window_test.py b/integrations/tensorflow/e2e/sliding_window_test.py index b663fc8..f206d86 100644 --- a/integrations/tensorflow/e2e/sliding_window_test.py +++ b/integrations/tensorflow/e2e/sliding_window_test.py
@@ -76,7 +76,7 @@ @tf_test_utils.compile_module(SlidingWindowModule, exported_names=["predict"]) -class SlidingWindowTest(tf_test_utils.SavedModelTestCase): +class SlidingWindowTest(tf_test_utils.CompiledModuleTestCase): def test_slidingwindow(self): input1 = np.array([[1.0, 2.0]], dtype=np.float32)
diff --git a/integrations/tensorflow/e2e/strings_test.py b/integrations/tensorflow/e2e/strings_test.py index ac590ff..ce0787e 100644 --- a/integrations/tensorflow/e2e/strings_test.py +++ b/integrations/tensorflow/e2e/strings_test.py
@@ -41,7 +41,7 @@ @tf_test_utils.compile_module(StringsModule) -class StringsTest(tf_test_utils.SavedModelTestCase): +class StringsTest(tf_test_utils.CompiledModuleTestCase): def test_print_ids(self): input_ids = np.asarray(
diff --git a/integrations/tensorflow/e2e/tensorlist_test.py b/integrations/tensorflow/e2e/tensorlist_test.py index f8ea811..9b1330c 100644 --- a/integrations/tensorflow/e2e/tensorlist_test.py +++ b/integrations/tensorflow/e2e/tensorlist_test.py
@@ -12,10 +12,6 @@ # See the License for the specific language governing permissions and # limitations under the License. -from __future__ import absolute_import -from __future__ import division -from __future__ import print_function - from pyiree.tf.support import tf_test_utils import tensorflow.compat.v2 as tf @@ -69,7 +65,7 @@ @tf_test_utils.compile_module(TensorListModule) -class TensorListTest(tf_test_utils.SavedModelTestCase): +class TensorListTest(tf_test_utils.CompiledModuleTestCase): def test_identity_through_tensorlist(self): m = self.get_module()
diff --git a/iree/base/BUILD b/iree/base/BUILD index b8c81ed..8d6b446 100644 --- a/iree/base/BUILD +++ b/iree/base/BUILD
@@ -43,6 +43,7 @@ ":file_mapping", ":init", ":tracing", + "@com_google_absl//absl/time", ], ) @@ -60,7 +61,6 @@ ":status", "@com_google_absl//absl/base:core_headers", "@com_google_absl//absl/container:inlined_vector", - "@com_google_absl//absl/time", ], ) @@ -431,7 +431,16 @@ name = "time", hdrs = ["time.h"], deps = [ - "@com_google_absl//absl/time", + ":api", + ], +) + +cc_test( + name = "time_test", + srcs = ["time_test.cc"], + deps = [ + ":time", + "//iree/testing:gtest_main", ], )
diff --git a/iree/base/CMakeLists.txt b/iree/base/CMakeLists.txt index 3a6ca54..322bf43 100644 --- a/iree/base/CMakeLists.txt +++ b/iree/base/CMakeLists.txt
@@ -38,6 +38,7 @@ ::file_mapping ::init ::tracing + absl::time PUBLIC ) @@ -519,10 +520,20 @@ HDRS "time.h" DEPS - absl::time + ::api PUBLIC ) +iree_cc_test( + NAME + time_test + SRCS + "time_test.cc" + DEPS + ::time + iree::testing::gtest_main +) + if(${IREE_ENABLE_RUNTIME_TRACING}) iree_cc_library( NAME
diff --git a/iree/base/api.cc b/iree/base/api.cc index ed74804..be992c4 100644 --- a/iree/base/api.cc +++ b/iree/base/api.cc
@@ -18,6 +18,7 @@ #include <cstring> #include <string> +#include "absl/time/clock.h" #include "iree/base/api_util.h" #include "iree/base/file_mapping.h" #include "iree/base/init.h" @@ -91,6 +92,24 @@ } //===----------------------------------------------------------------------===// +// iree_time_t and iree_duration_t +//===----------------------------------------------------------------------===// + +IREE_API_EXPORT iree_time_t iree_time_now() { + return absl::GetCurrentTimeNanos(); +} + +IREE_API_EXPORT iree_time_t +iree_relative_timeout_to_deadline_ns(iree_duration_t timeout_ns) { + if (timeout_ns == IREE_DURATION_ZERO) { + return IREE_TIME_INFINITE_PAST; + } else if (timeout_ns == IREE_DURATION_INFINITE) { + return IREE_TIME_INFINITE_FUTURE; + } + return iree_time_now() + timeout_ns; +} + +//===----------------------------------------------------------------------===// // iree_allocator_t //===----------------------------------------------------------------------===//
diff --git a/iree/base/api.h b/iree/base/api.h index 3bfcdf3..ba8ec9b 100644 --- a/iree/base/api.h +++ b/iree/base/api.h
@@ -167,21 +167,6 @@ // Whole length of the underlying buffer. #define IREE_WHOLE_BUFFER (iree_device_size_t(-1)) -// Like absl::Time, represented as nanoseconds since unix epoch. -// TODO(benvanik): pick something easy to get into/outof time_t/etc. -typedef int64_t iree_time_t; -// Like absl::InfinitePast. -#define IREE_TIME_INFINITE_PAST INT64_MIN -// Like absl::InfiniteFuture. -#define IREE_TIME_INFINITE_FUTURE INT64_MAX - -// Like absl::Duration, represented as relative nanoseconds. -typedef int64_t iree_duration_t; -// Like absl::InfiniteDuration. -#define IREE_DURATION_INFINITE INT64_MIN -// Like absl::ZeroDuration. -#define IREE_DURATION_ZERO 0 - // A span of mutable bytes (ala std::span of uint8_t). typedef struct { uint8_t* data; @@ -375,6 +360,43 @@ #endif // IREE_API_NO_PROTOTYPES //===----------------------------------------------------------------------===// +// iree_time_t and iree_duration_t +//===----------------------------------------------------------------------===// + +// Like absl::Time, represented as nanoseconds since unix epoch. +// TODO(benvanik): pick something easy to get into/outof time_t/etc. +typedef int64_t iree_time_t; +// Like absl::InfinitePast. +#define IREE_TIME_INFINITE_PAST INT64_MIN +// Like absl::InfiniteFuture. +#define IREE_TIME_INFINITE_FUTURE INT64_MAX + +// Like absl::Duration, represented as relative nanoseconds. +typedef int64_t iree_duration_t; +// Like absl::InfiniteDuration. +#define IREE_DURATION_INFINITE INT64_MAX +// Like absl::ZeroDuration. +#define IREE_DURATION_ZERO 0 + +#ifndef IREE_API_NO_PROTOTYPES + +// Returns the current system time in unix nanoseconds. +// Depending on the system architecture and power mode this time may have a +// very coarse granularity (on the order of microseconds to milliseconds). +// +// The system timer may not be monotonic; users should ensure when comparing +// times they check for negative values in case the time moves backwards. +IREE_API_EXPORT iree_time_t iree_time_now(); + +// Converts a relative timeout duration to an absolute deadline time. +// This handles the special cases of IREE_DURATION_ZERO and +// IREE_DURATION_INFINITE to avoid extraneous time queries. +IREE_API_EXPORT iree_time_t +iree_relative_timeout_to_deadline_ns(iree_duration_t timeout_ns); + +#endif // IREE_API_NO_PROTOTYPES + +//===----------------------------------------------------------------------===// // iree_allocator_t (std::allocator-like interface) //===----------------------------------------------------------------------===//
diff --git a/iree/base/api_util.h b/iree/base/api_util.h index 279a6fa..5ec86a5 100644 --- a/iree/base/api_util.h +++ b/iree/base/api_util.h
@@ -17,7 +17,6 @@ #include "absl/base/macros.h" #include "absl/container/inlined_vector.h" -#include "absl/time/time.h" #include "iree/base/api.h" #include "iree/base/logging.h" #include "iree/base/status.h" @@ -100,39 +99,6 @@ } \ lhs = std::move(statusor).value() -// Converts an iree_duration_t to its equivalent absl::Duration. -inline absl::Duration ToAbslDuration(iree_duration_t duration) { - if (duration == IREE_DURATION_ZERO) { - return absl::ZeroDuration(); - } else if (duration == IREE_DURATION_INFINITE) { - return absl::InfiniteDuration(); - } else { - return absl::Nanoseconds(duration); - } -} - -// Converts an iree_time_t to its equivalent absl::Time. -inline absl::Time ToAbslTime(iree_time_t time) { - if (time == IREE_TIME_INFINITE_PAST) { - return absl::InfinitePast(); - } else if (time == IREE_TIME_INFINITE_FUTURE) { - return absl::InfiniteFuture(); - } else { - return absl::FromUnixNanos(time); - } -} - -// Converts an absl::Time to an iree_time_t. -inline iree_time_t FromAbslTime(absl::Time time) { - if (time == absl::InfinitePast()) { - return IREE_TIME_INFINITE_PAST; - } else if (time == absl::InfiniteFuture()) { - return IREE_TIME_INFINITE_FUTURE; - } else { - return absl::ToUnixNanos(time); - } -} - // Returns a vector initialized with the contents of a C-style list query. // For functions of the form (..., capacity, out_values, out_count) this will // try to fetch the items and resize as needed such that the returned value
diff --git a/iree/base/atomics.h b/iree/base/atomics.h index 0bd99c7..47587dc 100644 --- a/iree/base/atomics.h +++ b/iree/base/atomics.h
@@ -39,7 +39,7 @@ #endif #if defined(IREE_COMPILER_CLANG) - +// Emulate C11 atomics with builtins. typedef _Atomic intptr_t iree_atomic_intptr_t; #define IREE_ATOMIC_VAR_INIT(value) (value) #define iree_atomic_load(object) __c11_atomic_load(object, __ATOMIC_SEQ_CST) @@ -51,7 +51,7 @@ __c11_atomic_fetch_sub(object, operand, __ATOMIC_SEQ_CST) #elif defined(IREE_COMPILER_MSVC) - +// Emulate C11 atomics with Interlocked win32 APIs. // NOTE: currently assumes sizeof(intptr_t) == 8. typedef struct { intptr_t __val; @@ -68,19 +68,27 @@ InterlockedExchangeAdd64((volatile LONGLONG*)object, -(operand)) #elif defined(IREE_COMPILER_GCC) - -typedef _Atomic __INTPTR_TYPE__ iree_atomic_intptr_t; +// Emulate atomics for GCC in a way that is compatible for inclusion in +// both C and C++ modes. +#ifdef __cplusplus +// Equiv to C++ auto keyword in C++ mode. +#define __iree_auto_type auto +#else +// Only defined in C mode. +#define __iree_auto_type __auto_type +#endif +typedef __INTPTR_TYPE__ iree_atomic_intptr_t; #define IREE_ATOMIC_VAR_INIT(value) (value) #define iree_atomic_load(object) \ __atomic_load_ptr(object, __ATOMIC_SEQ_CST) __extension__({ \ - __auto_type __atomic_load_ptr = (object); \ + __iree_auto_type __atomic_load_ptr = (object); \ __typeof__(*__atomic_load_ptr) __atomic_load_tmp; \ __atomic_load(__atomic_load_ptr, &__atomic_load_tmp, (__ATOMIC_SEQ_CST)); \ __atomic_load_tmp; \ }) #define iree_atomic_store(object, desired) \ __extension__({ \ - __auto_type __atomic_store_ptr = (object); \ + __iree_auto_type __atomic_store_ptr = (object); \ __typeof__(*__atomic_store_ptr) __atomic_store_tmp = (desired); \ __atomic_store(__atomic_store_ptr, &__atomic_store_tmp, \ (__ATOMIC_SEQ_CST)); \
diff --git a/iree/base/signature_mangle.h b/iree/base/signature_mangle.h index 948dbb5..8afb4b5 100644 --- a/iree/base/signature_mangle.h +++ b/iree/base/signature_mangle.h
@@ -156,7 +156,7 @@ // ----------------------------------------------------------------------------- // Mangles raw function signatures. -// See function_abi.md. +// See docs/design_docs/function_abi.md. class RawSignatureMangler { public: static SignatureBuilder ToFunctionSignature(const SignatureBuilder& inputs, @@ -364,7 +364,8 @@ // Mangles function signatures according to the Sip (Structured Index Path) V1 // scheme. // -// Mangler for the 'sip' ABI. See function_abi.md in the documentation. +// Mangler for the 'sip' ABI. See docs/design_docs/function_abi.md in the +// documentation. class SipSignatureMangler { public: enum class IndexMode { @@ -443,7 +444,8 @@ // Parser for signatures generated by SipSignatureMangler. // This uses a Visitor interface to walk either input or result structs. // -// Mangler for the 'sip' ABI. See function_abi.md in the documentation. +// Mangler for the 'sip' ABI. See docs/design_docs/function_abi.md in the +// documentation. class SipSignatureParser { public: enum class StructType {
diff --git a/iree/base/time.h b/iree/base/time.h index 81f9809..e51a5b1 100644 --- a/iree/base/time.h +++ b/iree/base/time.h
@@ -15,31 +15,143 @@ #ifndef IREE_BASE_TIME_H_ #define IREE_BASE_TIME_H_ -#include <chrono> // NOLINT -#include <thread> // NOLINT +#include <type_traits> +#include <utility> -#include "absl/time/clock.h" -#include "absl/time/time.h" +#include "iree/base/api.h" namespace iree { +namespace impl { +template <class Tag, typename T> +class ChronoType { + public: + ChronoType() : value_() {} + explicit ChronoType(const T& value) : value_(value) {} + explicit ChronoType(T&& value) noexcept( + std::is_nothrow_move_constructible<T>::value) + : value_(std::move(value)) {} -// Converts a relative timeout duration to an absolute deadline time. -// This handles the special cases of absl::ZeroDuration and -// absl::InfiniteDuration to avoid extraneous time queries. -inline absl::Time RelativeTimeoutToDeadline(absl::Duration timeout) { - if (timeout == absl::ZeroDuration()) { - return absl::InfinitePast(); - } else if (timeout == absl::InfiniteDuration()) { - return absl::InfiniteFuture(); + explicit operator T&() noexcept { return value_; } + explicit operator const T&() const noexcept { return value_; } + + friend void swap(ChronoType& a, ChronoType& b) noexcept { + using std::swap; + swap(static_cast<T&>(a), static_cast<T&>(b)); } - return absl::Now() + timeout; + + friend inline bool operator==(const ChronoType& lhs, const ChronoType& rhs) { + return lhs.value_ == rhs.value_; + } + friend inline bool operator!=(const ChronoType& lhs, const ChronoType& rhs) { + return !(lhs == rhs); + } + friend inline bool operator<(const ChronoType& lhs, const ChronoType& rhs) { + return lhs.value_ < rhs.value_; + } + friend inline bool operator>(const ChronoType& lhs, const ChronoType& rhs) { + return rhs < lhs; + } + friend inline bool operator<=(const ChronoType& lhs, const ChronoType& rhs) { + return !(lhs > rhs); + } + friend inline bool operator>=(const ChronoType& lhs, const ChronoType& rhs) { + return !(lhs < rhs); + } + + friend ChronoType& operator+=(ChronoType& lhs, const ChronoType& rhs) { + static_cast<T&>(lhs) += static_cast<const T&>(rhs); + return lhs; + } + friend ChronoType operator+(const ChronoType& lhs, const ChronoType& rhs) { + return ChronoType(static_cast<const T&>(lhs) + static_cast<const T&>(rhs)); + } + + friend ChronoType& operator-=(ChronoType& lhs, const ChronoType& rhs) { + static_cast<T&>(lhs) -= static_cast<const T&>(rhs); + return lhs; + } + friend ChronoType operator-(const ChronoType& lhs, const ChronoType& rhs) { + return ChronoType(static_cast<const T&>(lhs) - static_cast<const T&>(rhs)); + } + + private: + T value_; +}; +} // namespace impl + +struct Duration : public impl::ChronoType<Duration, iree_duration_t> { + using ChronoType::ChronoType; + explicit operator uint64_t() const noexcept { + if (static_cast<iree_duration_t>(*this) == IREE_DURATION_INFINITE) { + return UINT64_MAX; + } + int64_t relative_ns = static_cast<int64_t>(*this); + return relative_ns <= 0 ? 0 : static_cast<uint64_t>(relative_ns); + } +}; + +static inline Duration InfiniteDuration() { + return Duration(IREE_DURATION_INFINITE); +} +static inline Duration ZeroDuration() { return Duration(IREE_DURATION_ZERO); } + +struct Time : public impl::ChronoType<Time, iree_time_t> { + using ChronoType::ChronoType; + friend Duration operator+(const Time& lhs, const Time& rhs) { + if (static_cast<iree_time_t>(lhs) == IREE_TIME_INFINITE_FUTURE || + static_cast<iree_time_t>(rhs) == IREE_TIME_INFINITE_FUTURE) { + return InfiniteDuration(); + } else if (static_cast<iree_time_t>(lhs) == IREE_TIME_INFINITE_PAST || + static_cast<iree_time_t>(rhs) == IREE_TIME_INFINITE_PAST) { + return ZeroDuration(); + } + return Duration(static_cast<const iree_time_t&>(lhs) + + static_cast<const iree_time_t&>(rhs)); + } + friend Duration operator-(const Time& lhs, const Time& rhs) { + if (static_cast<iree_time_t>(lhs) == IREE_TIME_INFINITE_FUTURE || + static_cast<iree_time_t>(rhs) == IREE_TIME_INFINITE_FUTURE) { + return InfiniteDuration(); + } else if (static_cast<iree_time_t>(lhs) == IREE_TIME_INFINITE_PAST || + static_cast<iree_time_t>(rhs) == IREE_TIME_INFINITE_PAST) { + return ZeroDuration(); + } + return Duration(static_cast<const iree_time_t&>(lhs) - + static_cast<const iree_time_t&>(rhs)); + } +}; + +static inline Time InfinitePast() { return Time(IREE_TIME_INFINITE_PAST); } +static inline Time InfiniteFuture() { return Time(IREE_TIME_INFINITE_FUTURE); } + +static inline Duration Milliseconds(int64_t millis) { + return Duration(millis * 1000000ull); } -// Suspends execution of the calling thread for the given |duration|. -// Depending on platform this may have an extremely coarse resolution (upwards -// of several to dozens of milliseconds). -inline void Sleep(absl::Duration duration) { - std::this_thread::sleep_for(absl::ToChronoMilliseconds(duration)); +// Returns the current system time in unix nanoseconds. +// Depending on the system architecture and power mode this time may have a +// very coarse granularity (on the order of microseconds to milliseconds). +// +// The system timer may not be monotonic; users should ensure when comparing +// times they check for negative values in case the time moves backwards. +static inline Time Now() { return Time(iree_time_now()); } + +// Converts a relative timeout duration to an absolute deadline time. +// This handles the special cases of IREE_DURATION_ZERO and +// IREE_DURATION_INFINITE to avoid extraneous time queries. +static inline Time RelativeTimeoutToDeadlineNanos(Duration timeout_ns) { + return Time(iree_relative_timeout_to_deadline_ns( + static_cast<iree_duration_t>(timeout_ns))); +} + +static inline Duration DeadlineToRelativeTimeoutNanos(Time deadline_ns) { + if (deadline_ns == InfiniteFuture()) { + return InfiniteDuration(); + } else if (deadline_ns == InfinitePast()) { + return ZeroDuration(); + } else { + return Duration(static_cast<uint64_t>(deadline_ns - Now())); + } } } // namespace iree
diff --git a/iree/base/time_test.cc b/iree/base/time_test.cc new file mode 100644 index 0000000..114cd4e --- /dev/null +++ b/iree/base/time_test.cc
@@ -0,0 +1,45 @@ +// Copyright 2020 Google LLC +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// https://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#include "iree/base/time.h" + +#include "iree/testing/gtest.h" + +namespace iree { +namespace { + +TEST(Time, DurationComparisons) { + EXPECT_TRUE(Milliseconds(123) == Milliseconds(123)); + EXPECT_FALSE(Milliseconds(123) == Milliseconds(456)); + EXPECT_FALSE(Milliseconds(123) != Milliseconds(123)); + EXPECT_TRUE(Milliseconds(123) != Milliseconds(456)); + + EXPECT_TRUE(Milliseconds(123) < Milliseconds(456)); + EXPECT_FALSE(Milliseconds(123) > Milliseconds(456)); + EXPECT_FALSE(Milliseconds(123) > Milliseconds(123)); + EXPECT_FALSE(Milliseconds(123) < Milliseconds(123)); + + EXPECT_TRUE(Milliseconds(123) <= Milliseconds(123)); + EXPECT_TRUE(Milliseconds(123) >= Milliseconds(123)); + EXPECT_TRUE(Milliseconds(123) <= Milliseconds(456)); + EXPECT_FALSE(Milliseconds(123) >= Milliseconds(456)); +} + +TEST(Time, DurationArithmetic) { + EXPECT_EQ(Milliseconds(150), Milliseconds(100) + Milliseconds(50)); + EXPECT_EQ(Milliseconds(50), Milliseconds(100) - Milliseconds(50)); +} + +} // namespace +} // namespace iree
diff --git a/iree/base/wait_handle.cc b/iree/base/wait_handle.cc index 2c67f31..9f00a33 100644 --- a/iree/base/wait_handle.cc +++ b/iree/base/wait_handle.cc
@@ -76,7 +76,7 @@ // ppoll is preferred as it has a much better timing mechanism; poll can have a // large slop on the deadline. // Documentation: https://linux.die.net/man/2/poll -StatusOr<int> SystemPoll(absl::Span<pollfd> poll_fds, absl::Time deadline) { +StatusOr<int> SystemPoll(absl::Span<pollfd> poll_fds, Time deadline_ns) { // Convert the deadline into a tmo_p struct for ppoll that controls whether // the call is blocking or non-blocking. Note that we must do this every // iteration of the loop as a previous ppoll may have taken some of the @@ -86,16 +86,16 @@ // http://man7.org/linux/man-pages/man2/poll.2.html timespec timeout_spec; timespec* tmo_p; - if (deadline == absl::InfinitePast()) { + if (deadline == InfinitePast()) { // 0 for non-blocking. timeout_spec = {0}; tmo_p = &timeout_spec; - } else if (deadline == absl::InfiniteFuture()) { + } else if (deadline == InfiniteFuture()) { // nullptr to ppoll() to block forever. tmo_p = nullptr; } else { // Wait only for as much time as we have before the deadline is exceeded. - absl::Duration remaining_time = deadline - absl::Now(); + absl::Duration remaining_time = deadline - Now(); if (remaining_time < absl::ZeroDuration()) { // Note: we likely have already bailed before getting here with a negative // duration. @@ -111,16 +111,16 @@ // poll(), present pretty much everywhere. // Documentation: https://linux.die.net/man/2/poll -StatusOr<int> SystemPoll(absl::Span<pollfd> poll_fds, absl::Time deadline) { +StatusOr<int> SystemPoll(absl::Span<pollfd> poll_fds, Time deadline_ns) { int timeout; - if (deadline == absl::InfinitePast()) { + if (deadline == InfinitePast()) { // Don't block. timeout = 0; - } else if (deadline == absl::InfiniteFuture()) { + } else if (deadline == InfiniteFuture()) { // Block forever. timeout = -1; } else { - absl::Duration remaining_time = deadline - absl::Now(); + absl::Duration remaining_time = deadline - Now(); if (remaining_time < absl::ZeroDuration()) { return DeadlineExceededErrorBuilder(IREE_LOC); } @@ -139,7 +139,7 @@ // The provided deadline will be observed if any of the wait handles needs to // block for acquiring an fd. StatusOr<absl::FixedArray<pollfd>> AcquireWaitHandles( - WaitHandle::WaitHandleSpan wait_handles, absl::Time deadline) { + WaitHandle::WaitHandleSpan wait_handles, Time deadline_ns) { absl::FixedArray<pollfd> poll_fds{wait_handles.size()}; for (int i = 0; i < wait_handles.size(); ++i) { poll_fds[i].events = POLLIN | POLLPRI | POLLERR | POLLHUP | POLLNVAL; @@ -160,7 +160,7 @@ poll_fds[i].fd = fd_info.second; // Abort if deadline exceeded. - if (deadline != absl::InfinitePast() && deadline < absl::Now()) { + if (deadline != InfinitePast() && deadline < Now()) { return DeadlineExceededErrorBuilder(IREE_LOC) << "Deadline exceeded acquiring for fds"; } @@ -204,7 +204,7 @@ // Performs a single poll on multiple fds and returns information about the // signaled fds, if any. Status MultiPoll(WaitHandle::WaitHandleSpan wait_handles, - absl::Span<pollfd> poll_fds, absl::Time deadline, + absl::Span<pollfd> poll_fds, Time deadline_ns, int* out_any_signaled_index, int* out_unsignaled_count) { *out_any_signaled_index = -1; *out_unsignaled_count = 0; @@ -293,7 +293,7 @@ public: std::string DebugString() const override { return "signal"; } StatusOr<std::pair<FdType, int>> AcquireFdForWait( - absl::Time deadline) override { + Time deadline_ns) override { return std::make_pair(FdType::kPermanent, kSignaledFd); } StatusOr<bool> TryResolveWakeOnFd(int fd) override { return true; } @@ -308,7 +308,7 @@ public: std::string DebugString() const override { return "fail"; } StatusOr<std::pair<FdType, int>> AcquireFdForWait( - absl::Time deadline) override { + Time deadline_ns) override { return InternalErrorBuilder(IREE_LOC) << "AlwaysFailingObject"; } StatusOr<bool> TryResolveWakeOnFd(int fd) override { @@ -320,7 +320,7 @@ } // static -Status WaitHandle::WaitAll(WaitHandleSpan wait_handles, absl::Time deadline) { +Status WaitHandle::WaitAll(WaitHandleSpan wait_handles, Time deadline_ns) { if (wait_handles.empty()) return OkStatus(); // Build the list of pollfds to wait on. @@ -332,7 +332,7 @@ int any_signaled_index = 0; RETURN_IF_ERROR(MultiPoll(wait_handles, absl::MakeSpan(poll_fds), deadline, &any_signaled_index, &unsignaled_count)); - } while (unsignaled_count > 0 && absl::Now() < deadline); + } while (unsignaled_count > 0 && Now() < deadline); if (unsignaled_count == 0) { // All waits resolved. @@ -345,7 +345,7 @@ // static StatusOr<bool> WaitHandle::TryWaitAll(WaitHandleSpan wait_handles) { - auto status = WaitAll(wait_handles, absl::InfinitePast()); + auto status = WaitAll(wait_handles, InfinitePast()); if (status.ok()) { return true; } else if (IsDeadlineExceeded(status)) { @@ -356,7 +356,7 @@ // static StatusOr<int> WaitHandle::WaitAny(WaitHandleSpan wait_handles, - absl::Time deadline) { + Time deadline_ns) { if (wait_handles.empty()) { return InvalidArgumentErrorBuilder(IREE_LOC) << "At least one wait handle is required for WaitAny"; @@ -379,7 +379,7 @@ // static StatusOr<int> WaitHandle::TryWaitAny(WaitHandleSpan wait_handles) { - auto status_or = WaitAny(wait_handles, absl::InfinitePast()); + auto status_or = WaitAny(wait_handles, InfinitePast()); return IsDeadlineExceeded(status_or.status()) ? -1 : status_or; } @@ -418,7 +418,7 @@ } StatusOr<bool> WaitHandle::TryWait() { - auto status = WaitAll({this}, absl::InfinitePast()); + auto status = WaitAll({this}, InfinitePast()); if (status.ok()) { return true; } else if (IsDeadlineExceeded(status)) {
diff --git a/iree/base/wait_handle.h b/iree/base/wait_handle.h index d051e51..7d9a4af 100644 --- a/iree/base/wait_handle.h +++ b/iree/base/wait_handle.h
@@ -20,8 +20,6 @@ #include <string> #include <utility> -#include "absl/time/clock.h" -#include "absl/time/time.h" #include "absl/types/span.h" #include "iree/base/ref_ptr.h" #include "iree/base/status.h" @@ -42,7 +40,7 @@ // } // private: // StatusOr<std::pair<FdType, int>> AcquireFdForWait( -// absl::Time deadline) override { +// Time deadline_ns) override { // // If blocking traditionally do so now and then return this: // return std::make_pair(FdType::kPermanent, kSignaledFd); // // Otherwise, see ManualResetEvent for an example using fds. @@ -112,14 +110,14 @@ // // In cases where the file descriptor may not be available the call may block // until either it is available or the |deadline| has elapsed. Use - // absl::InfinitePast() to prevent blocking. + // InfinitePast() to prevent blocking. // // Returns a valid file descriptor or kInvalidFd as an indication that the // object should not be waited on (already signaled, etc). Can return // kSignaledFd to indicate that it's already known that the handle has been // signaled and the caller should resolve as if it caused a wake normally. virtual StatusOr<std::pair<FdType, int>> AcquireFdForWait( - absl::Time deadline) = 0; + Time deadline_ns) = 0; // Tries to resolve the object with the given |fd|. // In many cases this will no-op, however some types may require additional @@ -159,12 +157,12 @@ // Returns DEADLINE_EXCEEDED if the |deadline| elapses without all handles // having been signaled. Note that a subset of the |wait_handles| may have // been signaled and each can be queried to see which one. - static Status WaitAll(WaitHandleSpan wait_handles, absl::Time deadline); - static Status WaitAll(WaitHandleSpan wait_handles, absl::Duration timeout) { - return WaitAll(wait_handles, RelativeTimeoutToDeadline(timeout)); + static Status WaitAll(WaitHandleSpan wait_handles, Time deadline_ns); + static Status WaitAll(WaitHandleSpan wait_handles, Duration timeout_ns) { + return WaitAll(wait_handles, RelativeTimeoutToDeadlineNanos(timeout_ns)); } static Status WaitAll(WaitHandleSpan wait_handles) { - return WaitAll(wait_handles, absl::InfiniteFuture()); + return WaitAll(wait_handles, InfiniteFuture()); } // Tries waiting on the handles and returns immediately if it would have @@ -184,14 +182,13 @@ // // Returns DEADLINE_EXCEEDED if the |deadline| elapses without any handles // having been signaled. + static StatusOr<int> WaitAny(WaitHandleSpan wait_handles, Time deadline_ns); static StatusOr<int> WaitAny(WaitHandleSpan wait_handles, - absl::Time deadline); - static StatusOr<int> WaitAny(WaitHandleSpan wait_handles, - absl::Duration timeout) { - return WaitAny(wait_handles, RelativeTimeoutToDeadline(timeout)); + Duration timeout_ns) { + return WaitAny(wait_handles, RelativeTimeoutToDeadlineNanos(timeout_ns)); } static StatusOr<int> WaitAny(WaitHandleSpan wait_handles) { - return WaitAny(wait_handles, absl::InfiniteFuture()); + return WaitAny(wait_handles, InfiniteFuture()); } // Tries waiting for at least one handle to complete and returns immediately @@ -242,11 +239,11 @@ // Returns success if the wait is successful and the |wait_handle| was // signaled. Returns DEADLINE_EXCEEDED if the timeout elapses without the // handle having been signaled. - Status Wait(absl::Time deadline) { return WaitAll({this}, deadline); } - Status Wait(absl::Duration timeout) { - return WaitAll({this}, RelativeTimeoutToDeadline(timeout)); + Status Wait(Time deadline_ns) { return WaitAll({this}, deadline); } + Status Wait(Duration timeout_ns) { + return WaitAll({this}, RelativeTimeoutToDeadlineNanos(timeout_ns)); } - Status Wait() { return WaitAll({this}, absl::InfiniteFuture()); } + Status Wait() { return WaitAll({this}, InfiniteFuture()); } // Tries waiting on the handle and returns immediately if it would have // waited. The caller will not be blocked even if the handle has not yet been @@ -304,8 +301,7 @@ void Initialize(); void Dispose(); - StatusOr<std::pair<FdType, int>> AcquireFdForWait( - absl::Time deadline) override { + StatusOr<std::pair<FdType, int>> AcquireFdForWait(Time deadline_ns) override { return std::make_pair(fd_type_, fd_); } StatusOr<bool> TryResolveWakeOnFd(int fd) override { return true; }
diff --git a/iree/base/wait_handle_test.cc b/iree/base/wait_handle_test.cc index ea4c7f1..a9f1ad1 100644 --- a/iree/base/wait_handle_test.cc +++ b/iree/base/wait_handle_test.cc
@@ -90,7 +90,7 @@ TEST(WaitHandleTest, SingleWait) { WaitHandle wh; ASSERT_OK(wh.Wait()); - ASSERT_OK(wh.Wait(absl::Now() + absl::Seconds(1))); + ASSERT_OK(wh.Wait(Now() + absl::Seconds(1))); ASSERT_OK(wh.Wait(absl::Seconds(1))); ASSERT_STATUSOR_TRUE(wh.TryWait()); } @@ -110,21 +110,21 @@ WaitHandle wh1 = fence1.OnSet(); // Poll; should return immediately with timeout. - ASSERT_TRUE(IsDeadlineExceeded( - WaitHandle::WaitAll({&wh0, &wh1}, absl::InfinitePast()))); + ASSERT_TRUE( + IsDeadlineExceeded(WaitHandle::WaitAll({&wh0, &wh1}, InfinitePast()))); // Notify fence1. ASSERT_OK(fence1.Set()); // Poll; should return immediately with timeout as fence1 is not signaled. - ASSERT_TRUE(IsDeadlineExceeded( - WaitHandle::WaitAll({&wh0, &wh1}, absl::InfinitePast()))); + ASSERT_TRUE( + IsDeadlineExceeded(WaitHandle::WaitAll({&wh0, &wh1}, InfinitePast()))); // Notify fence0. ASSERT_OK(fence0.Set()); // Poll again; should return immediately with success. - ASSERT_OK(WaitHandle::WaitAll({&wh0, &wh1}, absl::InfinitePast())); + ASSERT_OK(WaitHandle::WaitAll({&wh0, &wh1}, InfinitePast())); } // Tests waiting when the first file handle is invalid. This is to verify a @@ -134,14 +134,14 @@ WaitHandle wh = fence.OnSet(); // Poll; should return immediately with timeout as fence is not signaled. - ASSERT_TRUE(IsDeadlineExceeded( - WaitHandle::WaitAll({nullptr, &wh}, absl::InfinitePast()))); + ASSERT_TRUE( + IsDeadlineExceeded(WaitHandle::WaitAll({nullptr, &wh}, InfinitePast()))); // Notify fence. ASSERT_OK(fence.Set()); // Poll again; should return immediately with success. - ASSERT_OK(WaitHandle::WaitAll({nullptr, &wh}, absl::InfinitePast())); + ASSERT_OK(WaitHandle::WaitAll({nullptr, &wh}, InfinitePast())); } // Tests exceeding the timeout deadline with WaitAll. @@ -151,25 +151,24 @@ // Wait with timeout on the unsignaled fence: // Via polling (should never block): - ASSERT_TRUE( - IsDeadlineExceeded(WaitHandle::WaitAll({&wh}, absl::InfinitePast()))); + ASSERT_TRUE(IsDeadlineExceeded(WaitHandle::WaitAll({&wh}, InfinitePast()))); ASSERT_STATUSOR_FALSE(WaitHandle::TryWaitAll({&wh})); // Via time in the near future (should block): ASSERT_TRUE( - IsDeadlineExceeded(WaitHandle::WaitAll({&wh}, absl::Milliseconds(250)))); + IsDeadlineExceeded(WaitHandle::WaitAll({&wh}, Milliseconds(250)))); // Via time in the past, should exceed deadline. ASSERT_TRUE( - IsDeadlineExceeded(WaitHandle::WaitAll({&wh}, absl::Milliseconds(-250)))); + IsDeadlineExceeded(WaitHandle::WaitAll({&wh}, Milliseconds(-250)))); // Notify and ensure no more timeouts. ASSERT_OK(fence.Set()); - ASSERT_OK(WaitHandle::WaitAll({&wh}, absl::InfinitePast())); + ASSERT_OK(WaitHandle::WaitAll({&wh}, InfinitePast())); ASSERT_STATUSOR_TRUE(WaitHandle::TryWaitAll({&wh})); - ASSERT_OK(WaitHandle::WaitAll({&wh}, absl::Milliseconds(250))); + ASSERT_OK(WaitHandle::WaitAll({&wh}, Milliseconds(250))); // Via time in the past, should exceed deadline even if signaled. ASSERT_TRUE( - IsDeadlineExceeded(WaitHandle::WaitAll({&wh}, absl::Milliseconds(-250)))); + IsDeadlineExceeded(WaitHandle::WaitAll({&wh}, Milliseconds(-250)))); } // Tests using WaitAll to wait on other threads. @@ -177,12 +176,12 @@ // Spin up two threads. ManualResetEvent fence0; std::thread t0{[&]() { - ::usleep(absl::ToInt64Microseconds(absl::Milliseconds(250))); + ::usleep(absl::ToInt64Microseconds(Milliseconds(250))); ASSERT_OK(fence0.Set()); }}; ManualResetEvent fence1; std::thread t1{[&]() { - ::usleep(absl::ToInt64Microseconds(absl::Milliseconds(250))); + ::usleep(absl::ToInt64Microseconds(Milliseconds(250))); ASSERT_OK(fence1.Set()); }}; @@ -200,8 +199,8 @@ ManualResetEvent fence; WaitHandle wh0 = fence.OnSet(); WaitHandle wh1 = fence.OnSet(); - ASSERT_TRUE(IsDeadlineExceeded( - WaitHandle::WaitAll({&wh0, &wh1}, absl::InfinitePast()))); + ASSERT_TRUE( + IsDeadlineExceeded(WaitHandle::WaitAll({&wh0, &wh1}, InfinitePast()))); ASSERT_OK(fence.Set()); ASSERT_OK(WaitHandle::WaitAll({&wh0, &wh1})); } @@ -210,8 +209,8 @@ TEST(WaitHandleTest, WaitAllSameHandle) { ManualResetEvent fence; WaitHandle wh = fence.OnSet(); - ASSERT_TRUE(IsDeadlineExceeded( - WaitHandle::WaitAll({&wh, &wh}, absl::InfinitePast()))); + ASSERT_TRUE( + IsDeadlineExceeded(WaitHandle::WaitAll({&wh, &wh}, InfinitePast()))); ASSERT_OK(fence.Set()); ASSERT_OK(WaitHandle::WaitAll({&wh, &wh})); } @@ -243,14 +242,14 @@ // Poll; should return immediately with timeout. ASSERT_TRUE(IsDeadlineExceeded( - WaitHandle::WaitAny({&wh0, &wh1}, absl::InfinitePast()).status())); + WaitHandle::WaitAny({&wh0, &wh1}, InfinitePast()).status())); // Notify fence1. ASSERT_OK(fence1.Set()); // Poll; should return immediately with fence1 signaled. ASSERT_OK_AND_ASSIGN(int index, - WaitHandle::WaitAny({&wh0, &wh1}, absl::InfinitePast())); + WaitHandle::WaitAny({&wh0, &wh1}, InfinitePast())); EXPECT_EQ(1, index); // Notify fence0. @@ -258,7 +257,7 @@ // Poll again; should return immediately; which one is signaled is undefined. ASSERT_OK_AND_ASSIGN(index, - WaitHandle::WaitAny({&wh0, &wh1}, absl::InfinitePast())); + WaitHandle::WaitAny({&wh0, &wh1}, InfinitePast())); ASSERT_TRUE(index == 0 || index == 1); } @@ -272,41 +271,39 @@ // Wait with timeout on the unsignaled fences: // Via polling (should never block): ASSERT_TRUE(IsDeadlineExceeded( - WaitHandle::WaitAny({&wh0, &wh1}, absl::InfinitePast()).status())); + WaitHandle::WaitAny({&wh0, &wh1}, InfinitePast()).status())); ASSERT_OK_AND_ASSIGN(int index, WaitHandle::TryWaitAny({&wh0, &wh1})); ASSERT_EQ(-1, index); // Via time in the near future (should block): ASSERT_TRUE(IsDeadlineExceeded( - WaitHandle::WaitAny({&wh0, &wh1}, absl::Milliseconds(250)).status())); + WaitHandle::WaitAny({&wh0, &wh1}, Milliseconds(250)).status())); // Notify one of the fences. Should return immediately. ASSERT_OK(fence1.Set()); ASSERT_OK_AND_ASSIGN(index, - WaitHandle::WaitAny({&wh0, &wh1}, absl::InfinitePast())); + WaitHandle::WaitAny({&wh0, &wh1}, InfinitePast())); ASSERT_EQ(1, index); ASSERT_OK_AND_ASSIGN(index, WaitHandle::TryWaitAny({&wh0, &wh1})); ASSERT_EQ(1, index); - ASSERT_OK_AND_ASSIGN( - index, WaitHandle::WaitAny({&wh0, &wh1}, absl::Milliseconds(250))); + ASSERT_OK_AND_ASSIGN(index, + WaitHandle::WaitAny({&wh0, &wh1}, Milliseconds(250))); ASSERT_EQ(1, index); // The unnotified fence should still timeout. - ASSERT_TRUE(IsDeadlineExceeded( - WaitHandle::WaitAny({&wh0}, absl::InfinitePast()).status())); + ASSERT_TRUE( + IsDeadlineExceeded(WaitHandle::WaitAny({&wh0}, InfinitePast()).status())); ASSERT_OK_AND_ASSIGN(index, WaitHandle::TryWaitAny({&wh0})); ASSERT_EQ(-1, index); ASSERT_TRUE(IsDeadlineExceeded( - WaitHandle::WaitAny({&wh0}, absl::Milliseconds(250)).status())); + WaitHandle::WaitAny({&wh0}, Milliseconds(250)).status())); // Notify last fence and ensure complete. ASSERT_OK(fence0.Set()); - ASSERT_OK_AND_ASSIGN(index, - WaitHandle::WaitAny({&wh0}, absl::InfinitePast())); + ASSERT_OK_AND_ASSIGN(index, WaitHandle::WaitAny({&wh0}, InfinitePast())); ASSERT_EQ(0, index); ASSERT_OK_AND_ASSIGN(index, WaitHandle::TryWaitAny({&wh0})); ASSERT_EQ(0, index); - ASSERT_OK_AND_ASSIGN(index, - WaitHandle::WaitAny({&wh0}, absl::Milliseconds(250))); + ASSERT_OK_AND_ASSIGN(index, WaitHandle::WaitAny({&wh0}, Milliseconds(250))); ASSERT_EQ(0, index); } @@ -316,13 +313,13 @@ // t1 will wait on t0 such that they will act in sequence. ManualResetEvent fence0; std::thread t0{[&]() { - ::usleep(absl::ToInt64Microseconds(absl::Milliseconds(250))); + ::usleep(absl::ToInt64Microseconds(Milliseconds(250))); ASSERT_OK(fence0.Set()); }}; ManualResetEvent fence1; std::thread t1{[&]() { ASSERT_OK(fence0.OnSet().Wait()); - ::usleep(absl::ToInt64Microseconds(absl::Milliseconds(250))); + ::usleep(absl::ToInt64Microseconds(Milliseconds(250))); ASSERT_OK(fence1.Set()); }}; @@ -346,7 +343,7 @@ WaitHandle wh0 = fence.OnSet(); WaitHandle wh1 = fence.OnSet(); ASSERT_TRUE(IsDeadlineExceeded( - WaitHandle::WaitAny({&wh0, &wh1}, absl::InfinitePast()).status())); + WaitHandle::WaitAny({&wh0, &wh1}, InfinitePast()).status())); ASSERT_OK(fence.Set()); ASSERT_OK_AND_ASSIGN(int index, WaitHandle::WaitAny({&wh0, &wh1})); ASSERT_TRUE(index == 0 || index == 1); @@ -357,7 +354,7 @@ ManualResetEvent fence; WaitHandle wh = fence.OnSet(); ASSERT_TRUE(IsDeadlineExceeded( - WaitHandle::WaitAny({&wh, &wh}, absl::InfinitePast()).status())); + WaitHandle::WaitAny({&wh, &wh}, InfinitePast()).status())); ASSERT_OK(fence.Set()); ASSERT_OK_AND_ASSIGN(int index, WaitHandle::WaitAny({&wh, &wh})); ASSERT_TRUE(index == 0 || index == 1); @@ -386,7 +383,7 @@ MOCK_METHOD(std::string, DebugString, (), (const, override)); MOCK_METHOD((StatusOr<std::pair<FdType, int>>), AcquireFdForWait, - (absl::Time deadline), (override)); + (Time deadline_ns), (override)); MOCK_METHOD(StatusOr<bool>, TryResolveWakeOnFd, (int fd), (override)); WaitHandle OnSomething() { return WaitHandle(add_ref(this)); } @@ -402,7 +399,7 @@ // Try waiting; we should see the AcquireFdForWait and then return because // the fd has not been resolved. - EXPECT_CALL(mwo, AcquireFdForWait(_)).WillOnce([&](absl::Time deadline) { + EXPECT_CALL(mwo, AcquireFdForWait(_)).WillOnce([&](Time deadline_ns) { // Return the valid FD from the MRE. return mre.AcquireFdForWait(deadline); }); @@ -413,7 +410,7 @@ // Try waiting again; we should get the AcquireFdForWait and then also get // the TryResolveWakeOnFd. - EXPECT_CALL(mwo, AcquireFdForWait(_)).WillOnce([&](absl::Time deadline) { + EXPECT_CALL(mwo, AcquireFdForWait(_)).WillOnce([&](Time deadline_ns) { // Return the valid (and now signaled) FD from the MRE. return mre.AcquireFdForWait(deadline); }); @@ -431,13 +428,13 @@ // Make the AcquireFdForWait take longer than the timeout. We should hit // deadline exceeded even though always_wait hasn't be signaled. - EXPECT_CALL(mwo, AcquireFdForWait(_)).WillOnce([](absl::Time deadline) { - ::usleep(absl::ToInt64Microseconds(absl::Milliseconds(10))); + EXPECT_CALL(mwo, AcquireFdForWait(_)).WillOnce([](Time deadline_ns) { + ::usleep(absl::ToInt64Microseconds(Milliseconds(10))); return std::make_pair(WaitableObject::FdType::kPermanent, WaitableObject::kInvalidFd); }); - ASSERT_TRUE(IsDeadlineExceeded(WaitHandle::WaitAll( - {&wh, &always_signal}, absl::Now() - absl::Milliseconds(250)))); + ASSERT_TRUE(IsDeadlineExceeded( + WaitHandle::WaitAll({&wh, &always_signal}, Now() - Milliseconds(250)))); } // Tests TryResolveWakeOnFd when a handle is a permanent kSignaledFd.
diff --git a/iree/compiler/Conversion/HLOToLinalg/HLOToLinalgOnBuffers.cpp b/iree/compiler/Conversion/HLOToLinalg/HLOToLinalgOnBuffers.cpp index 99f7eb1..90178fc 100644 --- a/iree/compiler/Conversion/HLOToLinalg/HLOToLinalgOnBuffers.cpp +++ b/iree/compiler/Conversion/HLOToLinalg/HLOToLinalgOnBuffers.cpp
@@ -27,6 +27,7 @@ #include "iree/compiler/Dialect/IREE/IR/IREEOps.h" #include "iree/compiler/Dialect/Shape/IR/ShapeOps.h" #include "llvm/ADT/APInt.h" +#include "llvm/ADT/STLExtras.h" #include "llvm/ADT/SetVector.h" #include "llvm/ADT/SmallVector.h" #include "mlir/Dialect/Linalg/IR/LinalgOps.h" @@ -444,7 +445,7 @@ rewriter.getI64IntegerAttr(1), // args_out rewriter.getArrayAttr(indexingMaps), getParallelAndReductionIterAttrs(rewriter, nloops, nonParallelLoops), - /*doc=*/nullptr, /*library_call=*/nullptr); + /*doc=*/nullptr, /*library_call=*/nullptr, /*symbol_source=*/nullptr); // Add a block to the region. auto *region = &linalgOp.region(); @@ -523,11 +524,6 @@ LogicalResult PadOpConversion::apply( mhlo::PadOp op, ArrayRef<Value> inputBuffers, ArrayRef<Value> resultBuffers, ConversionPatternRewriter &rewriter) const { - if (llvm::any_of(op.interior_padding().getValues<IntegerAttr>(), - [](auto attr) { return attr.getInt() != 0; })) - return op.emitError( - "pad op with non-zero interiror_padding is not supported"); - mhlo::PadOp::Adaptor adaptor(inputBuffers); auto loc = op.getLoc(); @@ -535,84 +531,31 @@ Value paddingVal = paddingConstVal ? rewriter.create<ConstantOp>(loc, paddingConstVal).getResult() - : adaptor.padding_value(); + : rewriter.create<LoadOp>(loc, adaptor.padding_value()); - auto operandType = adaptor.operand().getType().cast<ShapedType>(); - int rank = operandType.getRank(); - - SmallVector<Attribute, 2> indexingMaps; - indexingMaps.emplace_back(getPadOpInputIndexingMap(op, rank, rewriter)); - if (!paddingConstVal) { - indexingMaps.emplace_back(AffineMapAttr::get( - AffineMap::get(rank, /*symbolCount=*/0, rewriter.getContext()))); - } - indexingMaps.emplace_back(AffineMapAttr::get( - AffineMap::getMultiDimIdentityMap(rank, rewriter.getContext()))); - - SmallVector<Type, 2> resultTypes = {}; - SmallVector<Value, 2> linalgOpArgs = {adaptor.operand()}; - if (!paddingConstVal) linalgOpArgs.push_back(adaptor.padding_value()); - linalgOpArgs.push_back(resultBuffers[0]); - auto linalgOp = rewriter.create<linalg::IndexedGenericOp>( - loc, resultTypes, linalgOpArgs, - rewriter.getI64IntegerAttr(linalgOpArgs.size() - 1), // args_in - rewriter.getI64IntegerAttr(1), // args_out - rewriter.getArrayAttr(indexingMaps), - getParallelAndReductionIterAttrs(rewriter, rank, /*nReduction=*/0), - /*doc=*/nullptr, /*library_call=*/nullptr); - - // Add a block to the region. - auto *region = &linalgOp.region(); - auto *block = rewriter.createBlock(region, region->end()); - SmallVector<Type, 4> bodyArgTypes; - bodyArgTypes.append(rank, rewriter.getIndexType()); - bodyArgTypes.append(linalgOpArgs.size(), operandType.getElementType()); - block->addArguments(bodyArgTypes); - rewriter.setInsertionPointToEnd(block); - - // If the `index` of the result at a particular dimension i, is d_i, check if - // - // (d_i >= edge_padding_low[i]) && - // (d_i < (edge_padding_low[i] + operand_shape[i])). - // - // If true, then use the value of the operand, otherwise use the padding - // value. const auto &edgePaddingLow = op.edge_padding_low(); - const auto &edgePaddingHigh = op.edge_padding_high(); - - Type indexType = rewriter.getIndexType(); - Value cond = nullptr; - auto applyAndOp = [&](Value val) { - cond = cond ? rewriter.create<AndOp>(loc, cond, val) : val; - }; - for (int i = 0; i < rank; ++i) { - Value dim = block->getArgument(i); - int64_t paddingLow = edgePaddingLow.getValue<IntegerAttr>(i).getInt(); - int64_t paddingHigh = edgePaddingHigh.getValue<IntegerAttr>(i).getInt(); - auto low = rewriter.create<ConstantOp>( - loc, indexType, rewriter.getIntegerAttr(indexType, paddingLow)); - - // d_i < (edge_padding_low[i] + operand_shape[i]) - if (paddingLow != 0 && paddingHigh != 0) { - auto operandExtent = rewriter.create<DimOp>(loc, adaptor.operand(), i); - auto bound = rewriter.create<AddIOp>(loc, low, operandExtent); - auto checkUb = - rewriter.create<CmpIOp>(loc, CmpIPredicate::slt, dim, bound); - applyAndOp(checkUb); - } - - if (paddingLow != 0) { - // d_i >= edge_padding_low[i] - auto checkLb = rewriter.create<CmpIOp>(loc, CmpIPredicate::sge, dim, low); - applyAndOp(checkLb); - } + const auto &interiorPadding = op.interior_padding(); + SmallVector<Value, 3> offsets, sizes, strides; + for (auto it : llvm::enumerate(llvm::zip(edgePaddingLow, interiorPadding))) { + Value startIndex = rewriter.create<ConstantIndexOp>( + loc, std::get<0>(it.value()).getZExtValue()); + offsets.push_back(startIndex); + Value size = rewriter.create<DimOp>(loc, resultBuffers[0], it.index()); + sizes.push_back(size); + Value stride = rewriter.create<ConstantIndexOp>( + loc, std::get<1>(it.value()).getZExtValue() + 1); + strides.push_back(stride); } - Value inputVal = block->getArgument(rank); - if (!paddingConstVal) paddingVal = block->getArgument(rank + 1); - Value result = - cond ? rewriter.create<SelectOp>(loc, cond, inputVal, paddingVal) - : inputVal; - rewriter.create<linalg::YieldOp>(loc, result); + + // TODO(hanchung): Move SubViewOp this down to before where it is used. + // The pass for splitting dispatch function for vulkan requires no other ops + // interleave with Linalg structured ops, so put the SubViewOp in the + // beginning. + auto subViewOp = rewriter.create<SubViewOp>(loc, resultBuffers[0], offsets, + sizes, strides); + rewriter.create<linalg::FillOp>(loc, resultBuffers[0], paddingVal); + rewriter.create<linalg::CopyOp>(loc, inputBuffers[0], subViewOp); + return success(); } @@ -691,8 +634,9 @@ int batch = op.batch_dims().getSExtValue(); auto indexShapeType = adaptor.index().getType().dyn_cast<ShapedType>(); int nIndices = indexShapeType.getRank(); - if (batch < 0) - return op.emitError("expected batch_dims is greater than or equal to zero"); + auto inputShapeType = adaptor.input().getType().dyn_cast<ShapedType>(); + if (axis < 0) axis += inputShapeType.getRank(); + if (batch < 0) batch += nIndices; Location loc = op.getLoc(); Value output = op.getResult(); @@ -715,7 +659,7 @@ rewriter.getI64IntegerAttr(1), // args_out rewriter.getArrayAttr(indexingMaps), getParallelAndReductionIterAttrs(rewriter, rank, /*nReduction=*/0), - /*doc=*/nullptr, /*library_call=*/nullptr); + /*doc=*/nullptr, /*library_call=*/nullptr, /*symbol_source=*/nullptr); // Add a block to the region. auto *region = &linalgOp.region(); @@ -1019,7 +963,7 @@ rewriter.getArrayAttr(indexingMaps), getParallelAndReductionIterAttrs(rewriter, nInputRank, reductionDims.size()), - /*doc=*/nullptr, /*library_call=*/nullptr); + /*doc=*/nullptr, /*library_call=*/nullptr, /*symbol_source=*/nullptr); linalgOp.region().takeBody(reduceOp.body()); { @@ -1095,7 +1039,8 @@ op.getLoc(), ArrayRef<Type>(), opArgs, op.args_in(), op.args_out(), op.indexing_maps(), op.iterator_types(), /*doc=*/nullptr, - /*library_call=*/nullptr); + /*library_call=*/nullptr, + /*symbol_source=*/nullptr); // Move the region from the replaced op into the new op. unsigned numTensorOperands = op.getNumOperands(); // indexed_generic op has arguments for each index. In the case of generic
diff --git a/iree/compiler/Conversion/HLOToLinalg/test/pad.mlir b/iree/compiler/Conversion/HLOToLinalg/test/pad.mlir index a167ef4..8d1d96d 100644 --- a/iree/compiler/Conversion/HLOToLinalg/test/pad.mlir +++ b/iree/compiler/Conversion/HLOToLinalg/test/pad.mlir
@@ -5,7 +5,8 @@ func @pad_cst() { %c0 = constant 0 : index %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<12x4xf32> - // CHECK: linalg.indexed_generic + // CHECK: linalg.fill + // CHECK: linalg.copy %1 = constant dense<0.0> : tensor<f32> %2 = "mhlo.pad"(%0, %1) { edge_padding_high = dense<[2, 3]> : tensor<2xi64>, @@ -29,7 +30,8 @@ %c0 = constant 0 : index %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<12x4xf32> %1 = hal.interface.load.tensor @legacy_io::@arg1, offset = %c0 : tensor<f32> - // CHECK: linalg.indexed_generic + // CHECK: linalg.fill + // CHECK: linalg.copy %2 = "mhlo.pad"(%0, %1) { edge_padding_high = dense<[2, 3]> : tensor<2xi64>, edge_padding_low = dense<[4, 5]> : tensor<2xi64>, @@ -52,7 +54,8 @@ func @pad_no_op() { %c0 = constant 0 : index %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<12x4xf32> - // CHECK: linalg.indexed_generic + // CHECK: linalg.fill + // CHECK: linalg.copy %1 = constant dense<0.0> : tensor<f32> %2 = "mhlo.pad"(%0, %1) { edge_padding_high = dense<0> : tensor<2xi64>,
diff --git a/iree/compiler/Conversion/LinalgToLLVM/BUILD b/iree/compiler/Conversion/LinalgToLLVM/BUILD index 41d9988..fcd9476 100644 --- a/iree/compiler/Conversion/LinalgToLLVM/BUILD +++ b/iree/compiler/Conversion/LinalgToLLVM/BUILD
@@ -22,7 +22,6 @@ name = "LinalgToLLVM", srcs = [ "ConvertToLLVM.cpp", - "HALInterfaceToMemrefArguments.cpp", "Passes.cpp", ], hdrs = [
diff --git a/iree/compiler/Conversion/LinalgToLLVM/CMakeLists.txt b/iree/compiler/Conversion/LinalgToLLVM/CMakeLists.txt index d21dc19..bc31e4e 100644 --- a/iree/compiler/Conversion/LinalgToLLVM/CMakeLists.txt +++ b/iree/compiler/Conversion/LinalgToLLVM/CMakeLists.txt
@@ -21,7 +21,6 @@ "Passes.h" SRCS "ConvertToLLVM.cpp" - "HALInterfaceToMemrefArguments.cpp" "Passes.cpp" DEPS MLIRAffineToStandard
diff --git a/iree/compiler/Conversion/LinalgToLLVM/ConvertToLLVM.cpp b/iree/compiler/Conversion/LinalgToLLVM/ConvertToLLVM.cpp index 193eb3c..8686dca 100644 --- a/iree/compiler/Conversion/LinalgToLLVM/ConvertToLLVM.cpp +++ b/iree/compiler/Conversion/LinalgToLLVM/ConvertToLLVM.cpp
@@ -124,6 +124,180 @@ } }; +/// Returns true if `aOp` has a desciptor (set, binding) pair smaller than +/// `bOp`. Note that this ignores the offset. +bool operator<(IREE::HAL::InterfaceBindingOp aOp, + IREE::HAL::InterfaceBindingOp bOp) { + if (aOp.set().getZExtValue() == bOp.set().getZExtValue()) + return aOp.binding().getZExtValue() < bOp.binding().getZExtValue(); + return aOp.set().getZExtValue() < bOp.set().getZExtValue(); +} + +// Change signature of entry function to func +// entry_func(%packed_buffers_arg_ptr: +// !<llvm.int8**>, %push_constant: !<llvm.int64*>) and lower IREE and HAL ops to +// corresponding LLVMIR ops to construct memref descriptors and load +// push_constant values. +class ConvertFuncWithHALInterface : public ConvertToLLVMPattern { + public: + explicit ConvertFuncWithHALInterface(MLIRContext *context, + LLVMTypeConverter &typeconverter) + : ConvertToLLVMPattern(FuncOp::getOperationName(), context, + typeconverter) {} + + LogicalResult matchAndRewrite( + Operation *op, ArrayRef<Value> operands, + ConversionPatternRewriter &rewriter) const override { + if (SymbolTable::getSymbolVisibility(op) != SymbolTable::Visibility::Public) + return failure(); + auto funcOp = dyn_cast_or_null<FuncOp>(op); + FunctionType fnType = funcOp.getType(); + if (fnType.getNumInputs() != 0) { + return rewriter.notifyMatchFailure( + funcOp, "entry function should not have inputs"); + } + + // Get interface buffers from all the blocks. + SmallVector<IREE::PlaceholderOp, 8> bufferOps; + SmallVector<IREE::HAL::InterfaceLoadConstantOp, 8> loadOps; + for (Block &block : funcOp.getBlocks()) { + for (Operation &op : block) { + if (auto phOp = dyn_cast<IREE::PlaceholderOp>(op)) + bufferOps.push_back(phOp); + if (auto phOp = dyn_cast<IREE::HAL::InterfaceLoadConstantOp>(op)) { + loadOps.push_back(phOp); + } + } + } + + if (bufferOps.empty()) return failure(); + + // A map from buffer ops to their corresponding interface binding ops. + llvm::DenseMap<Operation *, IREE::HAL::InterfaceBindingOp> bufferBindingMap; + for (auto bufferOp : bufferOps) { + auto symbol = SymbolTable::lookupNearestSymbolFrom( + bufferOp, bufferOp.getAttrOfType<SymbolRefAttr>("binding")); + bufferBindingMap[bufferOp] = cast<IREE::HAL::InterfaceBindingOp>(symbol); + } + + // Sort buffers according to their descriptor (set, binding) pair. + llvm::sort(bufferOps, [&bufferBindingMap](IREE::PlaceholderOp aBuffer, + IREE::PlaceholderOp bBuffer) { + return bufferBindingMap[aBuffer] < bufferBindingMap[bBuffer]; + }); + + // A map from buffer ops to their corresponding function argument indices. + llvm::DenseMap<Operation *, unsigned> bufferArgMap; + // A map from binding ops to their corresponding function argument indices. + llvm::DenseMap<Operation *, unsigned> bindingArgMap; + llvm::SmallVector<MemRefType, 4> inputMemRefTypes; + llvm::SmallVector<LLVM::LLVMType, 4> inputStructPtrs; + unsigned argIndex = 0; + for (auto bufferOp : bufferOps) { + auto binding = bufferBindingMap[bufferOp]; + auto it = bindingArgMap.find(binding); + if (it != bindingArgMap.end()) { + bufferArgMap[bufferOp] = it->second; + } else { + bindingArgMap[binding] = argIndex; + bufferArgMap[bufferOp] = argIndex; + ++argIndex; + } + + auto memrefType = bufferOp.getType().dyn_cast_or_null<MemRefType>(); + inputMemRefTypes.push_back(memrefType); + auto elementType = typeConverter.convertType(memrefType.getElementType()) + .dyn_cast<LLVM::LLVMType>(); + if (!elementType) return failure(); + inputStructPtrs.push_back( + elementType.getPointerTo(memrefType.getMemorySpace())); + } + + TypeConverter::SignatureConversion signatureConverter(/*numOrigInputs=*/0); + + // func foo(%packed_buffer_args: !llvm<i8**>, %push_constant: !llvm<i64*>) + auto packedBuffersArgsTy = + LLVM::LLVMType::getInt8PtrTy(typeConverter.getDialect()).getPointerTo(); + auto pushConstantArgTy = + LLVM::LLVMType::getInt64Ty(typeConverter.getDialect()).getPointerTo(); + signatureConverter.addInputs(packedBuffersArgsTy); + signatureConverter.addInputs(pushConstantArgTy); + + // Create the new function's signature. + Location loc = funcOp.getLoc(); + auto newFuncOp = rewriter.create<FuncOp>( + loc, funcOp.getName(), + rewriter.getFunctionType(signatureConverter.getConvertedTypes(), + llvm::None), + ArrayRef<NamedAttribute>()); + + // Move all ops in the old function's region to the new function. + rewriter.inlineRegionBefore(funcOp.getBody(), newFuncOp.getBody(), + newFuncOp.end()); + rewriter.applySignatureConversion(&newFuncOp.getBody(), signatureConverter); + + auto builder = OpBuilder::atBlockBegin(&(newFuncOp.getBlocks().front())); + + // Cast and unpack input packed_buffer_arguments and construct memref + // descriptors. + Value packedBuffersArgsPtr = builder.create<LLVM::BitcastOp>( + loc, + LLVM::LLVMType::getStructTy(typeConverter.getDialect(), inputStructPtrs) + .getPointerTo(), + newFuncOp.getArgument(0)); + Value packedBuffersArgs = + builder.create<LLVM::LoadOp>(loc, packedBuffersArgsPtr); + for (auto bufferOp : bufferOps) { + MemRefType memrefType = bufferOp.getType().dyn_cast_or_null<MemRefType>(); + if (!memrefType) return failure(); + const auto index = bufferArgMap[bufferOp]; + Value bufferPtr = builder.create<LLVM::ExtractValueOp>( + loc, inputStructPtrs[index], packedBuffersArgs, + rewriter.getI64ArrayAttr(index)); + if (memrefType.hasStaticShape()) { + auto desc = MemRefDescriptor::fromStaticShape( + builder, loc, typeConverter, memrefType, bufferPtr); + rewriter.replaceOp(bufferOp, {desc}); + } else { + auto desc = MemRefDescriptor::undef( + builder, loc, typeConverter.convertType(memrefType)); + desc.setAllocatedPtr(builder, loc, bufferPtr); + desc.setAlignedPtr(builder, loc, bufferPtr); + rewriter.replaceOp(bufferOp, {desc}); + } + } + + // Lower hal.interface.load.constant ops into llvm.getelementptr, llvm.load + for (auto loadOp : loadOps) { + Value offset = builder.create<LLVM::ConstantOp>( + loc, LLVM::LLVMType::getInt64Ty(typeConverter.getDialect()), + builder.getI64IntegerAttr(loadOp.offset().getZExtValue())); + Value constPtr = builder.create<LLVM::GEPOp>(loc, pushConstantArgTy, + newFuncOp.getArgument(1), + ArrayRef<Value>({offset})); + Value dimConstant = builder.create<LLVM::LoadOp>(loc, constPtr); + rewriter.replaceOp(loadOp, dimConstant); + } + + rewriter.eraseOp(funcOp); + return success(); + } +}; + +class RemoveInterfaceOpPattern : public ConvertToLLVMPattern { + public: + explicit RemoveInterfaceOpPattern(MLIRContext *context, + LLVMTypeConverter &typeconverter) + : ConvertToLLVMPattern(IREE::HAL::InterfaceOp::getOperationName(), + context, typeconverter) {} + LogicalResult matchAndRewrite( + Operation *op, ArrayRef<Value> operands, + ConversionPatternRewriter &rewriter) const override { + rewriter.eraseOp(op); + return success(); + } +}; + namespace { struct ConvertToLLVMPass : public PassWrapper<ConvertToLLVMPass, OperationPass<ModuleOp>> { @@ -145,11 +319,12 @@ populateVectorToLLVMConversionPatterns(converter, patterns); populateLinalgToLLVMConversionPatterns(converter, patterns, &getContext()); // The following patterns resolves dynamic shapes by substituting tie_shape - // ops with an updated memref descriptors and replacing RankDimOp with actual - // index loaded from memref<?xi32> that holds all dynamic shapes - // push constants. - patterns.insert<ConvertRankedDimPattern, ConvertTieShapePattern, - RemoveMakeRankedShape>(&getContext(), converter); + // ops with an updated memref descriptors and replacing RankDimOp with + // actual index loaded from memref<?xi32> that holds all dynamic shapes push + // constants. + patterns.insert<ConvertFuncWithHALInterface, ConvertRankedDimPattern, + ConvertTieShapePattern, RemoveMakeRankedShape, + RemoveInterfaceOpPattern>(&getContext(), converter); LLVMConversionTarget target(getContext()); target.addLegalOp<ModuleOp, ModuleTerminatorOp>(); if (failed(applyPartialConversion(module, target, patterns))) @@ -162,7 +337,8 @@ static PassRegistration<ConvertToLLVMPass> pass( "iree-codegen-convert-to-llvm", - "Perform final conversion from Linalg/HAL/Shape/Vector/Standard to LLVMIR " + "Perform final conversion from Linalg/HAL/Shape/Vector/Standard to " + "LLVMIR " "dialect", [] { return std::make_unique<ConvertToLLVMPass>(); });
diff --git a/iree/compiler/Conversion/LinalgToLLVM/HALInterfaceToMemrefArguments.cpp b/iree/compiler/Conversion/LinalgToLLVM/HALInterfaceToMemrefArguments.cpp deleted file mode 100644 index ac968e4..0000000 --- a/iree/compiler/Conversion/LinalgToLLVM/HALInterfaceToMemrefArguments.cpp +++ /dev/null
@@ -1,231 +0,0 @@ -// Copyright 2020 Google LLC -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// https://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. - -#include <memory> - -#include "iree/compiler/Dialect/HAL/IR/HALDialect.h" -#include "iree/compiler/Dialect/HAL/IR/HALOps.h" -#include "iree/compiler/Dialect/IREE/IR/IREEOps.h" -#include "mlir/Dialect/StandardOps/IR/Ops.h" -#include "mlir/IR/Builders.h" -#include "mlir/Pass/Pass.h" -#include "mlir/Pass/PassRegistry.h" -#include "mlir/Transforms/DialectConversion.h" - -namespace mlir { -namespace iree_compiler { -namespace { - -/// Returns true if the given function contains interface related operations -/// that are used by other ops. -bool containsUsedInterfaceOp(FuncOp funcOp) { - for (Block& block : funcOp.getBlocks()) { - for (Operation& op : block) { - if (!op.getUses().empty() && - (isa<IREE::PlaceholderOp>(op) || - isa<IREE::HAL::InterfaceLoadConstantOp>(op))) { - return true; - } - } - } - return false; -} - -/// Returns true if `aOp` has a desciptor (set, binding) pair smaller than -/// `bOp`. Note that this ignores the offset. -bool operator<(IREE::HAL::InterfaceBindingOp aOp, - IREE::HAL::InterfaceBindingOp bOp) { - if (aOp.set().getZExtValue() == bOp.set().getZExtValue()) - return aOp.binding().getZExtValue() < bOp.binding().getZExtValue(); - return aOp.set().getZExtValue() < bOp.set().getZExtValue(); -} - -/// A pattern to process function interface. It replaces interface related ops -/// with function arguments to match LLVM's CodeGen's ABI contract. -/// -/// IREE scheduler passes interface ABI information via hal.interface.* ops to -/// all backends. We create iree.placeholder ops to represent buffers behind -/// those hal.interface.* ops. However the LLVM CodeGen uses function parameters -/// and memref descriptors for ABI. So we need to bridge the gap somewhere. -/// -/// This pass finds all interface buffers used in the function, sort them -/// according to the descriptor (set, binding) pair, and put unique ones as -/// function parameters in order. -/// Note: This should be kept consistent with LLVM's HAL backend. -struct ProcessFuncInterfacePattern : public OpConversionPattern<FuncOp> { - using OpConversionPattern::OpConversionPattern; - LogicalResult matchAndRewrite( - FuncOp funcOp, ArrayRef<Value> Operands, - ConversionPatternRewriter& rewriter) const override { - // Only process entry functions. - if (SymbolTable::getSymbolVisibility(funcOp) != - SymbolTable::Visibility::Public) - return failure(); - - FunctionType fnType = funcOp.getType(); - if (fnType.getNumInputs() != 0) - return rewriter.notifyMatchFailure( - funcOp, "entry function should not have inputs"); - - // Get interface buffers from all the blocks. - SmallVector<IREE::PlaceholderOp, 8> bufferOps; - SmallVector<IREE::HAL::InterfaceLoadConstantOp, 8> loadOps; - for (Block& block : funcOp.getBlocks()) { - for (Operation& op : block) { - if (auto phOp = dyn_cast<IREE::PlaceholderOp>(op)) - bufferOps.push_back(phOp); - if (auto phOp = dyn_cast<IREE::HAL::InterfaceLoadConstantOp>(op)) { - loadOps.push_back(phOp); - } - } - } - - if (bufferOps.empty()) return failure(); - - // A map from buffer ops to their corresponding interface binding ops. - llvm::DenseMap<Operation*, IREE::HAL::InterfaceBindingOp> bufferBindingMap; - for (auto bufferOp : bufferOps) { - auto symbol = SymbolTable::lookupNearestSymbolFrom( - bufferOp, bufferOp.getAttrOfType<SymbolRefAttr>("binding")); - bufferBindingMap[bufferOp] = cast<IREE::HAL::InterfaceBindingOp>(symbol); - } - - // Sort buffers according to their descriptor (set, binding) pair. - llvm::sort(bufferOps, [&bufferBindingMap](IREE::PlaceholderOp aBuffer, - IREE::PlaceholderOp bBuffer) { - return bufferBindingMap[aBuffer] < bufferBindingMap[bBuffer]; - }); - - // Create a function argument for each of the unique binding pointed by the - // buffer ops. - TypeConverter::SignatureConversion signatureConverter(/*numOrigInputs=*/0); - // A map from buffer ops to their corresponding function argument indices. - llvm::DenseMap<Operation*, unsigned> bufferArgMap; - // A map from binding ops to their corresponding function argument indices. - llvm::DenseMap<Operation*, unsigned> bindingArgMap; - unsigned argIndex = 0; - for (auto bufferOp : bufferOps) { - auto binding = bufferBindingMap[bufferOp]; - auto it = bindingArgMap.find(binding); - if (it != bindingArgMap.end()) { - bufferArgMap[bufferOp] = it->second; - } else { - bindingArgMap[binding] = argIndex; - bufferArgMap[bufferOp] = argIndex; - signatureConverter.addInputs(bufferOp.getType()); - ++argIndex; - } - } - Type dynamicDimsBufferType = - MemRefType::get(ShapedType::kDynamicSize, rewriter.getIntegerType(32)); - signatureConverter.addInputs(dynamicDimsBufferType); - - // Create the new function's signature. - Location loc = funcOp.getLoc(); - auto newFuncOp = rewriter.create<FuncOp>( - loc, funcOp.getName(), - rewriter.getFunctionType(signatureConverter.getConvertedTypes(), - llvm::None), - ArrayRef<NamedAttribute>()); - newFuncOp.setAttr("llvm.emit_c_interface", - mlir::UnitAttr::get(funcOp.getContext())); - - // Move all ops in the old function's region to the new function. - rewriter.inlineRegionBefore(funcOp.getBody(), newFuncOp.getBody(), - newFuncOp.end()); - rewriter.applySignatureConversion(&newFuncOp.getBody(), signatureConverter); - - // Replace all buffer ops' uses with the newly created function arguments - // and erase them. - for (auto bufferOp : bufferOps) { - bufferOp.replaceAllUsesWith( - newFuncOp.getArgument(bufferArgMap[bufferOp])); - - rewriter.eraseOp(bufferOp); - } - - // Lower all hal.interface.load.constant ops into std.load - // from the last buffer holding all dynamic dimensions with the proper - // offset. - Type indexType = rewriter.getIndexType(); - auto builder = OpBuilder::atBlockBegin(&(newFuncOp.getBlocks().front())); - auto newLoc = newFuncOp.front().front().getLoc(); - for (auto loadOp : loadOps) { - SmallVector<Value, 1> indices; - Value constantOffset = builder.create<ConstantOp>( - newLoc, indexType, - rewriter.getIntegerAttr(indexType, loadOp.offset().getZExtValue())); - indices.push_back(constantOffset); - Value loadDim = builder.create<LoadOp>( - newLoc, newFuncOp.getArgument(newFuncOp.getNumArguments() - 1), - indices); - Value loadDimIndex = - builder.create<IndexCastOp>(newLoc, loadDim, indexType); - loadOp.replaceAllUsesWith(loadDimIndex); - rewriter.eraseOp(loadOp); - } - rewriter.eraseOp(funcOp); - return success(); - } -}; - -struct RemoveInterfaceOpPattern - : public OpRewritePattern<IREE::HAL::InterfaceOp> { - using OpRewritePattern::OpRewritePattern; - LogicalResult matchAndRewrite(IREE::HAL::InterfaceOp interfaceOp, - PatternRewriter& rewriter) const override { - rewriter.eraseOp(interfaceOp); - return success(); - } -}; - -/// Converting from Linalg to LLVM needs to run on a module and since it -/// applies a full conversion, make a module with jst the impl function. -struct HALInterfaceToMemrefArgumentsPass - : PassWrapper<HALInterfaceToMemrefArgumentsPass, OperationPass<ModuleOp>> { - void runOnOperation() override { - MLIRContext& context = getContext(); - - OwningRewritePatternList patterns; - patterns.insert<ProcessFuncInterfacePattern>(&context); - patterns.insert<RemoveInterfaceOpPattern>(&context); - - ConversionTarget target(context); - // Convert the interface related ops away. - target.addDynamicallyLegalOp<FuncOp>( - [](FuncOp funcOp) { return !containsUsedInterfaceOp(funcOp); }); - target.addIllegalOp<IREE::PlaceholderOp>(); - target.addIllegalDialect<IREE::HAL::HALDialect>(); - // Allow the rest. - target.markUnknownOpDynamicallyLegal([](Operation*) { return true; }); - - if (failed(applyFullConversion(getOperation(), target, patterns))) - return signalPassFailure(); - } -}; - -} // namespace - -std::unique_ptr<OperationPass<ModuleOp>> -createHALInterfaceToMemrefArgumentsPass() { - return std::make_unique<HALInterfaceToMemrefArgumentsPass>(); -} - -static PassRegistration<HALInterfaceToMemrefArgumentsPass> pass( - "iree-codegen-hal-interface-to-memref-arguments-pass", - "Convert a function with HAL bindings interface to memref arguments", - [] { return std::make_unique<HALInterfaceToMemrefArgumentsPass>(); }); - -} // namespace iree_compiler -} // namespace mlir
diff --git a/iree/compiler/Conversion/LinalgToLLVM/Passes.cpp b/iree/compiler/Conversion/LinalgToLLVM/Passes.cpp index e8c6d9c..8c8eb21 100644 --- a/iree/compiler/Conversion/LinalgToLLVM/Passes.cpp +++ b/iree/compiler/Conversion/LinalgToLLVM/Passes.cpp
@@ -35,10 +35,7 @@ passManager.addPass(createCanonicalizerPass()); passManager.addPass(createCSEPass()); - // Convert ExecuableOp entry function to use memref arguments. - passManager.addPass(createHALInterfaceToMemrefArgumentsPass()); - - // (Linalg, STD) -> LLVM + // (HAL, IREE, Linalg, STD) -> LLVM // OpPassManager& llvmPassManager = passManager.nest<ModuleOp>(); passManager.addPass(createConvertToLLVMPass()); passManager.addPass(createCanonicalizerPass());
diff --git a/iree/compiler/Conversion/LinalgToLLVM/Passes.h b/iree/compiler/Conversion/LinalgToLLVM/Passes.h index 5bfb893..fdad0e6 100644 --- a/iree/compiler/Conversion/LinalgToLLVM/Passes.h +++ b/iree/compiler/Conversion/LinalgToLLVM/Passes.h
@@ -20,11 +20,6 @@ namespace mlir { namespace iree_compiler { -/// Converts function signture type from hal interface op annotation to memref -/// argument. -std::unique_ptr<OperationPass<ModuleOp>> -createHALInterfaceToMemrefArgumentsPass(); - /// Pass to perform final conversion to LLVM dialect. std::unique_ptr<OperationPass<ModuleOp>> createConvertToLLVMPass();
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/ConvertToGPUPass.cpp b/iree/compiler/Conversion/LinalgToSPIRV/ConvertToGPUPass.cpp index c0d5135..81514ec 100644 --- a/iree/compiler/Conversion/LinalgToSPIRV/ConvertToGPUPass.cpp +++ b/iree/compiler/Conversion/LinalgToSPIRV/ConvertToGPUPass.cpp
@@ -17,6 +17,9 @@ // Partition computation within dispatch function to workgroups/workitems. // //===----------------------------------------------------------------------===// + +#include <array> + #include "iree/compiler/Conversion/LinalgToSPIRV/Attributes.h" #include "iree/compiler/Conversion/LinalgToSPIRV/MarkerUtils.h" #include "iree/compiler/Conversion/LinalgToSPIRV/Passes.h" @@ -425,7 +428,7 @@ unsigned numDims, MutableArrayRef<Value> id, MutableArrayRef<Value> count) { - ArrayRef<StringRef> dims = {"x", "y", "z"}; + std::array<StringRef, 3> dims{"x", "y", "z"}; assert(id.size() == numDims); assert(count.size() == numDims); for (unsigned i = 0; i < numDims; ++i) { @@ -560,7 +563,7 @@ ConversionPatternRewriter &rewriter) const override { // Check for marker that specifies that the linalg op is to be partitioned // across threads within a workgroup. - if (!hasWorkItemMarker(linalgOp)) return failure(); + if (!hasWorkGroupMarker(linalgOp)) return failure(); Optional<linalg::LinalgLoops> loops = linalg::linalgLowerOpToLoops<scf::ParallelOp>(rewriter, linalgOp); if (!loops) return failure(); @@ -584,7 +587,7 @@ LogicalResult matchAndRewrite( LinalgOpTy linalgOp, ArrayRef<Value> operands, ConversionPatternRewriter &rewriter) const override { - if (!hasWorkItemMarker(linalgOp)) return failure(); + if (!hasWorkGroupMarker(linalgOp)) return failure(); Optional<linalg::LinalgLoops> loops = linalg::linalgLowerOpToLoops<scf::ParallelOp>(rewriter, linalgOp); if (!loops) return failure();
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/LinalgTileAndFusePass.cpp b/iree/compiler/Conversion/LinalgToSPIRV/LinalgTileAndFusePass.cpp index e9dddd6..934e5ae 100644 --- a/iree/compiler/Conversion/LinalgToSPIRV/LinalgTileAndFusePass.cpp +++ b/iree/compiler/Conversion/LinalgToSPIRV/LinalgTileAndFusePass.cpp
@@ -314,7 +314,7 @@ LogicalResult matchAndRewrite(Operation *op, PatternRewriter &rewriter) const override { - if (!hasWorkItemMarker(op)) return failure(); + if (!hasWorkGroupMarker(op)) return failure(); return linalg::LinalgPromotionPattern<linalg::MatmulOp>::matchAndRewrite( op, rewriter); } @@ -365,7 +365,7 @@ .setLoopType(linalg::LinalgTilingLoopType::ParallelLoops), tileSizeCalculator.getWorkGroupSize(), linalg::LinalgMarker(ArrayRef<Identifier>(), - Identifier::get(getWorkItemMarker(), context))); + Identifier::get(getWorkGroupMarker(), context))); applyPatternsAndFoldGreedily(getOperation(), tilingPatterns); if (useWorkgroupMemory) { @@ -385,7 +385,7 @@ [&](OpBuilder &b, Value src, Value dst) -> LogicalResult { return copyToFromWorkgroupMemory(b, src, dst); }), - linalg::LinalgMarker(Identifier::get(getWorkItemMarker(), context), + linalg::LinalgMarker(Identifier::get(getWorkGroupMarker(), context), Identifier::get(PromotionMarker, context))); applyPatternsAndFoldGreedily(getOperation(), promotionPatterns); } @@ -394,7 +394,7 @@ OpBuilder builder(context); funcOp.walk([&builder](linalg::LinalgOp linalgOp) { if (hasMarker(linalgOp, PromotionMarker)) { - setWorkItemMarker(linalgOp); + setWorkGroupMarker(linalgOp); insertBarrierAfter(builder, linalgOp.getLoc(), linalgOp); } });
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/MarkerUtils.cpp b/iree/compiler/Conversion/LinalgToSPIRV/MarkerUtils.cpp index c874234..47747de 100644 --- a/iree/compiler/Conversion/LinalgToSPIRV/MarkerUtils.cpp +++ b/iree/compiler/Conversion/LinalgToSPIRV/MarkerUtils.cpp
@@ -36,8 +36,6 @@ StringRef getWorkGroupMarker() { return "workgroup"; } -StringRef getWorkItemMarker() { return "workitem"; } - bool hasMarker(Operation *op, StringRef marker) { return checkMarkerValue(op, marker); } @@ -46,10 +44,6 @@ return checkMarkerValue(op, getWorkGroupMarker()); } -bool hasWorkItemMarker(Operation *op) { - return checkMarkerValue(op, getWorkItemMarker()); -} - void setMarker(Operation *op, StringRef marker) { op->setAttr(linalg::LinalgTransforms::kLinalgTransformMarker, StringAttr::get(marker, op->getContext())); @@ -57,6 +51,5 @@ void setWorkGroupMarker(Operation *op) { setMarker(op, getWorkGroupMarker()); } -void setWorkItemMarker(Operation *op) { setMarker(op, getWorkItemMarker()); } } // namespace iree_compiler } // namespace mlir
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/MarkerUtils.h b/iree/compiler/Conversion/LinalgToSPIRV/MarkerUtils.h index 36dccca..e512ead 100644 --- a/iree/compiler/Conversion/LinalgToSPIRV/MarkerUtils.h +++ b/iree/compiler/Conversion/LinalgToSPIRV/MarkerUtils.h
@@ -31,7 +31,7 @@ namespace iree_compiler { /// Marker to denote that a linalg operation is to be partitioned to workitems. -StringRef getWorkItemMarker(); +StringRef getWorkGroupMarker(); /// Returns true if an operation has the specified `marker`. When `marker` is /// empty, returns true if the operation has any marker. @@ -39,14 +39,14 @@ /// Returns true if an operation has marker to denote that it is to be /// partitioned to workitems. -bool hasWorkItemMarker(Operation *); +bool hasWorkGroupMarker(Operation *); /// Sets a given marker on an operation. void setMarker(Operation *, StringRef); /// Sets marker to denote that a linalg operation is to be partitioned to /// workitems. -void setWorkItemMarker(Operation *); +void setWorkGroupMarker(Operation *); } // namespace iree_compiler } // namespace mlir
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/convert_to_gpu.mlir b/iree/compiler/Conversion/LinalgToSPIRV/test/convert_to_gpu.mlir index 679f523..64621f3 100644 --- a/iree/compiler/Conversion/LinalgToSPIRV/test/convert_to_gpu.mlir +++ b/iree/compiler/Conversion/LinalgToSPIRV/test/convert_to_gpu.mlir
@@ -162,7 +162,7 @@ %12 = dim %arg2, %c1 : memref<?x?xf32> %13 = affine.min #map0(%arg4)[%12] %14 = subview %arg2[%arg3, %arg4] [%11, %13] [1, 1] : memref<?x?xf32> to memref<?x?xf32, #map2> - linalg.matmul %5, %9, %14 {__internal_linalg_transform__ = "workitem"} : (memref<?x?xf32, #map2>, memref<?x?xf32, #map2>, memref<?x?xf32, #map2>) + linalg.matmul %5, %9, %14 {__internal_linalg_transform__ = "workgroup"} : (memref<?x?xf32, #map2>, memref<?x?xf32, #map2>, memref<?x?xf32, #map2>) } scf.yield } @@ -235,7 +235,7 @@ %13 = affine.min #map5(%arg5)[%4] %14 = dim %arg2, %c3 : memref<?x?x?x?xf32> %15 = subview %arg2[%arg3, %arg4, %arg5, 0] [%11, %12, %13, %14] [1, 1, 1, 1] : memref<?x?x?x?xf32> to memref<?x?x?x?xf32, #map3> - linalg.conv(%arg0, %9, %15) {__internal_linalg_transform__ = "workitem", dilations = [1, 1], strides = [1, 1]} : memref<?x?x?x?xf32>, memref<?x?x?x?xf32, #map3>, memref<?x?x?x?xf32, #map3> + linalg.conv(%arg0, %9, %15) {__internal_linalg_transform__ = "workgroup", dilations = [1, 1], strides = [1, 1]} : memref<?x?x?x?xf32>, memref<?x?x?x?xf32, #map3>, memref<?x?x?x?xf32, #map3> scf.yield } return @@ -364,7 +364,7 @@ %9 = affine.min #map3(%arg3)[%2] %10 = affine.min #map4(%arg4)[%3] %11 = subview %arg2[%arg3, %arg4] [%9, %10] [1, 1] : memref<?x?xf32> to memref<?x?xf32, #map2> - linalg.pooling_max(%8, %arg1, %11) {__internal_linalg_transform__ = "workitem", dilations = [1, 1], strides = [1, 1]} : memref<?x?xf32, #map2>, memref<?x?xf32>, memref<?x?xf32, #map2> + linalg.pooling_max(%8, %arg1, %11) {__internal_linalg_transform__ = "workgroup", dilations = [1, 1], strides = [1, 1]} : memref<?x?xf32, #map2>, memref<?x?xf32>, memref<?x?xf32, #map2> scf.yield } return
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/convert_to_gpu_option.mlir b/iree/compiler/Conversion/LinalgToSPIRV/test/convert_to_gpu_option.mlir index 1701535..63f8aa5 100644 --- a/iree/compiler/Conversion/LinalgToSPIRV/test/convert_to_gpu_option.mlir +++ b/iree/compiler/Conversion/LinalgToSPIRV/test/convert_to_gpu_option.mlir
@@ -32,7 +32,7 @@ %13 = affine.min #map5(%arg5)[%4] %14 = dim %arg2, %c3 : memref<?x?x?x?xf32> %15 = subview %arg2[%arg3, %arg4, %arg5, 0] [%11, %12, %13, %14] [1, 1, 1, 1] : memref<?x?x?x?xf32> to memref<?x?x?x?xf32, #map3> - linalg.conv(%arg0, %9, %15) {__internal_linalg_transform__ = "workitem", dilations = [1, 1], strides = [1, 1]} : memref<?x?x?x?xf32>, memref<?x?x?x?xf32, #map3>, memref<?x?x?x?xf32, #map3> + linalg.conv(%arg0, %9, %15) {__internal_linalg_transform__ = "workgroup", dilations = [1, 1], strides = [1, 1]} : memref<?x?x?x?xf32>, memref<?x?x?x?xf32, #map3>, memref<?x?x?x?xf32, #map3> scf.yield } return
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/cyclic_to_workgroup.mlir b/iree/compiler/Conversion/LinalgToSPIRV/test/cyclic_to_workgroup.mlir index 110ac24..cac18ab 100644 --- a/iree/compiler/Conversion/LinalgToSPIRV/test/cyclic_to_workgroup.mlir +++ b/iree/compiler/Conversion/LinalgToSPIRV/test/cyclic_to_workgroup.mlir
@@ -27,7 +27,7 @@ %12 = dim %arg2, %c1 : memref<?x?xf32> %13 = affine.min #map0(%arg4)[%12] %14 = subview %arg2[%arg3, %arg4] [%11, %13] [1, 1] : memref<?x?xf32> to memref<?x?xf32, #map2> - linalg.matmul %5, %9, %14 {__internal_linalg_transform__ = "workitem"} : (memref<?x?xf32, #map2>, memref<?x?xf32, #map2>, memref<?x?xf32, #map2>) + linalg.matmul %5, %9, %14 {__internal_linalg_transform__ = "workgroup"} : (memref<?x?xf32, #map2>, memref<?x?xf32, #map2>, memref<?x?xf32, #map2>) } scf.yield }
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/linalg_tile_and_fuse.mlir b/iree/compiler/Conversion/LinalgToSPIRV/test/linalg_tile_and_fuse.mlir index 0e2fe6d..1728d35 100644 --- a/iree/compiler/Conversion/LinalgToSPIRV/test/linalg_tile_and_fuse.mlir +++ b/iree/compiler/Conversion/LinalgToSPIRV/test/linalg_tile_and_fuse.mlir
@@ -51,7 +51,7 @@ // CHECK: %[[VIEW2:.+]] = subview %[[ARG2]] // CHECK: linalg.conv // CHECK-SAME: %[[ARG0]], %[[VIEW1]], %[[VIEW2]] -// CHECK-SAME: "workitem" +// CHECK-SAME: "workgroup" // ----- @@ -81,7 +81,7 @@ // CHECK: %[[VIEW1:.+]] = subview %[[ARG1]] // CHECK: %[[VIEW2:.+]] = subview %[[ARG2]] // CHECK: linalg.matmul -// CHECK-SAME: "workitem" +// CHECK-SAME: "workgroup" // CHECK-SAME: %[[VIEW0]], %[[VIEW1]], %[[VIEW2]] // ----- @@ -111,4 +111,4 @@ // CHECK: %[[VIEW2:.+]] = subview %[[ARG2]] // CHECK: linalg.pooling_max // CHECK-SAME: %[[VIEW0]], %[[ARG1]], %[[VIEW2]] -// CHECK-SAME: "workitem" +// CHECK-SAME: "workgroup"
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/workgroup_memory_promotion.mlir b/iree/compiler/Conversion/LinalgToSPIRV/test/workgroup_memory_promotion.mlir index 76cfcb8..a24c77b 100644 --- a/iree/compiler/Conversion/LinalgToSPIRV/test/workgroup_memory_promotion.mlir +++ b/iree/compiler/Conversion/LinalgToSPIRV/test/workgroup_memory_promotion.mlir
@@ -36,12 +36,12 @@ // CHECK: %[[ALLOC2:.+]] = alloc(%[[C4]], %[[C8]]) : memref<?x?xf32, 3> // CHECK: %[[SUBVIEW2:.+]] = subview %[[ALLOC2]] // CHECK: linalg.copy(%[[ARG0SV]], %[[SUBVIEW1]]) -// CHECK-SAME: "workitem" +// CHECK-SAME: "workgroup" // CHECK: spv.ControlBarrier "Workgroup", "Workgroup", "AcquireRelease" // CHECK: linalg.copy(%[[ARG1SV]], %[[SUBVIEW2]]) -// CHECK-SAME: "workitem" +// CHECK-SAME: "workgroup" // CHECK: spv.ControlBarrier "Workgroup", "Workgroup", "AcquireRelease" -// CHECK: linalg.matmul {{.*}}"workitem"{{.*}} %[[SUBVIEW1]], %[[SUBVIEW2]], %[[RET0SV]] +// CHECK: linalg.matmul {{.*}}"workgroup"{{.*}} %[[SUBVIEW1]], %[[SUBVIEW2]], %[[RET0SV]] // CHECK: spv.ControlBarrier "Workgroup", "Workgroup", "AcquireRelease" // CHECK-DAG: dealloc %[[ALLOC1]] : memref<?x?xf32, 3> // CHECK-DAG: dealloc %[[ALLOC2]] : memref<?x?xf32, 3>
diff --git a/iree/compiler/Conversion/init_conversions.h b/iree/compiler/Conversion/init_conversions.h index 7a190e7..259e3d5 100644 --- a/iree/compiler/Conversion/init_conversions.h +++ b/iree/compiler/Conversion/init_conversions.h
@@ -47,7 +47,6 @@ inline void registerLinalgToLLVMPasses() { static bool init_once = []() { // LinalgToLLVM - createHALInterfaceToMemrefArgumentsPass(); return true; }(); (void)init_once;
diff --git a/iree/compiler/Dialect/Flow/IR/FlowOps.td b/iree/compiler/Dialect/Flow/IR/FlowOps.td index 23bce9c..13ad2f1 100644 --- a/iree/compiler/Dialect/Flow/IR/FlowOps.td +++ b/iree/compiler/Dialect/Flow/IR/FlowOps.td
@@ -610,6 +610,17 @@ let hasFolder = 1; } +def FLOW_TensorTraceOp : FLOW_Op<"tensor.trace", []> { + let summary = [{trace value(s) operation}]; + let description = [{ + Trace point for dispatchable functions. + }]; + + let arguments = (ins Variadic<FLOW_Tensor>:$operands); + + let assemblyFormat = "attr-dict ($operands^ `:` type($operands))?"; +} + //===----------------------------------------------------------------------===// // Streams //===----------------------------------------------------------------------===//
diff --git a/iree/compiler/Dialect/Flow/Transforms/OutlineDispatchRegions.cpp b/iree/compiler/Dialect/Flow/Transforms/OutlineDispatchRegions.cpp index 44a19d1..3281c6c 100644 --- a/iree/compiler/Dialect/Flow/Transforms/OutlineDispatchRegions.cpp +++ b/iree/compiler/Dialect/Flow/Transforms/OutlineDispatchRegions.cpp
@@ -25,6 +25,7 @@ #include "mlir/IR/BlockAndValueMapping.h" #include "mlir/IR/Builders.h" #include "mlir/IR/Diagnostics.h" +#include "mlir/IR/StandardTypes.h" #include "mlir/IR/SymbolTable.h" #include "mlir/Pass/Pass.h" @@ -40,6 +41,11 @@ namespace { +static llvm::cl::opt<bool> traceDispatchTensors( + "iree-flow-trace-dispatch-tensors", + llvm::cl::desc("Trace input/output values for each dispatch function"), + llvm::cl::init(false)); + // Converts a dispatch_region into a dispatch to the outlined region function. LogicalResult convertToDispatchOp(DispatchRegionOp regionOp, ExecutableOp executableOp, @@ -57,11 +63,28 @@ return failure(); } + auto getTensorTypeArgs = [](auto args) { + SmallVector<Value, 4> res; + for (auto arg : args) { + if (arg.getType().template isa<TensorType>()) res.push_back(arg); + } + return res; + }; + if (traceDispatchTensors) { + builder.create<TensorTraceOp>(regionOp.getLoc(), + getTensorTypeArgs(newArgs)); + } + // Create the dispatch op to the executable function. auto dispatchOp = builder.create<DispatchOp>( regionOp.getLoc(), executableOp.getName(), entryPointOp.getName(), regionOp.workload(), outlinedFuncOp.getType().getResults(), newArgs); + if (traceDispatchTensors) { + builder.create<TensorTraceOp>(regionOp.getLoc(), + getTensorTypeArgs(dispatchOp.getResults())); + } + // Replace uses of the existing results with the new results. for (int i = 0; i < regionOp.getNumResults(); ++i) { regionOp.getResult(i).replaceAllUsesWith(dispatchOp.getResult(i));
diff --git a/iree/compiler/Dialect/HAL/Conversion/FlowToHAL/ConvertTensorOps.cpp b/iree/compiler/Dialect/HAL/Conversion/FlowToHAL/ConvertTensorOps.cpp index 2947b70..6fcadbd 100644 --- a/iree/compiler/Dialect/HAL/Conversion/FlowToHAL/ConvertTensorOps.cpp +++ b/iree/compiler/Dialect/HAL/Conversion/FlowToHAL/ConvertTensorOps.cpp
@@ -19,6 +19,7 @@ #include "iree/compiler/Dialect/HAL/Utils/TypeUtils.h" #include "iree/compiler/Dialect/IREE/IR/IREETypes.h" #include "llvm/ADT/DenseMap.h" +#include "llvm/ADT/STLExtras.h" #include "mlir/Dialect/StandardOps/IR/Ops.h" #include "mlir/IR/Attributes.h" #include "mlir/IR/BlockAndValueMapping.h" @@ -120,13 +121,36 @@ } }; +class TensorTraceOpConversion + : public OpConversionPattern<IREE::Flow::TensorTraceOp> { + public: + TensorTraceOpConversion(MLIRContext *ctx, TypeConverter &converter) + : OpConversionPattern(ctx) {} + + LogicalResult matchAndRewrite( + IREE::Flow::TensorTraceOp traceOp, llvm::ArrayRef<Value> rawOperands, + ConversionPatternRewriter &rewriter) const override { + Location loc = traceOp.getLoc(); + SmallVector<Value, 4> bufferViews; + for (auto operand : llvm::enumerate(rawOperands)) { + auto adaptor = IREE::HAL::TensorRewriteAdaptor::get( + loc, traceOp.getOperand(operand.index()), operand.value(), rewriter); + bufferViews.emplace_back(adaptor.getBufferView()); + } + rewriter.replaceOpWithNewOp<IREE::HAL::BufferViewTraceOp>(traceOp, + bufferViews); + return success(); + } +}; + } // namespace void populateFlowTensorToHALPatterns(MLIRContext *context, OwningRewritePatternList &patterns, TypeConverter &converter) { patterns.insert<ConstantTensorOpConversion, TensorLoadOpConversion, - TensorStoreOpConversion>(context, converter); + TensorStoreOpConversion, TensorTraceOpConversion>(context, + converter); } } // namespace iree_compiler
diff --git a/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertBufferViewOps.cpp b/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertBufferViewOps.cpp index 7212549..5abbfa4 100644 --- a/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertBufferViewOps.cpp +++ b/iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertBufferViewOps.cpp
@@ -118,6 +118,8 @@ context, importSymbols, typeConverter, "hal.buffer_view.dim"); patterns.insert<BufferViewDimsOpConversion>( context, importSymbols, typeConverter, "hal.buffer_view.dims"); + patterns.insert<VMImportOpConversion<IREE::HAL::BufferViewTraceOp>>( + context, importSymbols, typeConverter, "hal.buffer_view.trace"); } } // namespace iree_compiler
diff --git a/iree/compiler/Dialect/HAL/IR/HALOps.td b/iree/compiler/Dialect/HAL/IR/HALOps.td index 0139ddb..9593647 100644 --- a/iree/compiler/Dialect/HAL/IR/HALOps.td +++ b/iree/compiler/Dialect/HAL/IR/HALOps.td
@@ -927,6 +927,17 @@ let assemblyFormat = [{$buffer_view attr-dict `:` type($result)}]; } +def HAL_BufferViewTraceOp : HAL_Op<"buffer_view.trace", []> { + let summary = [{trace value(s) operation}]; + let description = [{ + Trace point for dispatchable functions. + }]; + + let arguments = (ins Variadic<HAL_BufferView>:$operands); + + let assemblyFormat = "attr-dict ($operands^ `:` type($operands))?"; +} + //===----------------------------------------------------------------------===// // iree::hal::CommandBuffer //===----------------------------------------------------------------------===//
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/BUILD b/iree/compiler/Dialect/HAL/Target/LLVM/BUILD index 2e322af..0f6be4a 100644 --- a/iree/compiler/Dialect/HAL/Target/LLVM/BUILD +++ b/iree/compiler/Dialect/HAL/Target/LLVM/BUILD
@@ -105,9 +105,7 @@ cc_library( name = "LLVMAOTTargetLinker", hdrs = ["LLVMAOTTargetLinker.h"], - deps = [ - "//iree/base:file_io", - ] + platform_trampoline_deps("LLVMAOTTargetLinker", "compiler/Dialect/HAL/Target/LLVM"), + deps = platform_trampoline_deps("LLVMAOTTargetLinker", "compiler/Dialect/HAL/Target/LLVM"), ) cc_library( @@ -115,6 +113,6 @@ hdrs = ["LLVMAOTTargetLinker.h"], deps = [ ":LLVMTargetOptions", - "//iree/base:file_io", + "//iree/base:status", ], )
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/CMakeLists.txt b/iree/compiler/Dialect/HAL/Target/LLVM/CMakeLists.txt index 4aa0ad3..0ee00e4 100644 --- a/iree/compiler/Dialect/HAL/Target/LLVM/CMakeLists.txt +++ b/iree/compiler/Dialect/HAL/Target/LLVM/CMakeLists.txt
@@ -97,7 +97,6 @@ HDRS "LLVMAOTTargetLinker.h" DEPS - iree::base::file_io iree::compiler::Dialect::HAL::Target::LLVM::internal::LLVMAOTTargetLinker_internal PUBLIC ) @@ -109,6 +108,6 @@ "LLVMAOTTargetLinker.h" DEPS ::LLVMTargetOptions - iree::base::file_io + iree::base::status PUBLIC )
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/LLVMAOTTarget.cpp b/iree/compiler/Dialect/HAL/Target/LLVM/LLVMAOTTarget.cpp index 7269089..8cb47b5 100644 --- a/iree/compiler/Dialect/HAL/Target/LLVM/LLVMAOTTarget.cpp +++ b/iree/compiler/Dialect/HAL/Target/LLVM/LLVMAOTTarget.cpp
@@ -65,13 +65,10 @@ auto executableOp = cast<ExecutableOp>(targetOp.getParentOp()); auto entryPointOps = executableOp.getBlock().getOps<ExecutableEntryPointOp>(); - const bool addCInterface = true; + for (auto entryPointOp : entryPointOps) { - std::string funcName = - addCInterface ? "_mlir_ciface_" + std::string(entryPointOp.sym_name()) - : std::string(entryPointOp.sym_name()); - dyLibExecutableDef.entry_points.push_back("invoke_" + funcName); - createLLVMInvocationFunc(funcName, llvmModule.get()); + dyLibExecutableDef.entry_points.push_back( + std::string(entryPointOp.sym_name())); } // LLVMIR opt passes.
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/LLVMAOTTargetLinker.h b/iree/compiler/Dialect/HAL/Target/LLVM/LLVMAOTTargetLinker.h index 764ad02..669f17c 100644 --- a/iree/compiler/Dialect/HAL/Target/LLVM/LLVMAOTTargetLinker.h +++ b/iree/compiler/Dialect/HAL/Target/LLVM/LLVMAOTTargetLinker.h
@@ -18,7 +18,7 @@ #include <string> -#include "iree/base/file_io.h" +#include "iree/base/status.h" #include "iree/compiler/Dialect/HAL/Target/LLVM/LLVMTargetOptions.h" namespace mlir {
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/LLVMIRPasses.cpp b/iree/compiler/Dialect/HAL/Target/LLVM/LLVMIRPasses.cpp index cb2a526..e91441d 100644 --- a/iree/compiler/Dialect/HAL/Target/LLVM/LLVMIRPasses.cpp +++ b/iree/compiler/Dialect/HAL/Target/LLVM/LLVMIRPasses.cpp
@@ -44,44 +44,6 @@ return machine; } -void createLLVMInvocationFunc(const std::string& name, llvm::Module* module) { - // TODO(ataei): This is written as a stub in LLVM IR. It would be easier to - // have this using MLIR and lower it to LLVM like the dispatch function - // implementation is. - - auto& ctx = module->getContext(); - llvm::IRBuilder<> builder(ctx); - auto var_func = module->getFunction(name); - - auto new_type = llvm::FunctionType::get( - builder.getVoidTy(), builder.getInt8PtrTy()->getPointerTo(), - /*isVarArg=*/false); - - auto new_name = "invoke_" + name; - auto func_cst = module->getOrInsertFunction(new_name, new_type); - llvm::Function* interface_func = - llvm::cast<llvm::Function>(func_cst.getCallee()); - - auto bb = llvm::BasicBlock::Create(ctx); - bb->insertInto(interface_func); - builder.SetInsertPoint(bb); - llvm::Value* argList = interface_func->arg_begin(); - llvm::SmallVector<llvm::Value*, 8> args; - args.reserve(llvm::size(var_func->args())); - for (auto& indexedArg : llvm::enumerate(var_func->args())) { - llvm::Value* arg_index = llvm::Constant::getIntegerValue( - builder.getInt64Ty(), llvm::APInt(64, indexedArg.index())); - llvm::Value* arg_ptr_ptr = builder.CreateGEP(argList, arg_index); - llvm::Value* arg_ptr = builder.CreateLoad(arg_ptr_ptr); - arg_ptr = builder.CreateBitCast( - arg_ptr, indexedArg.value().getType()->getPointerTo()); - llvm::Value* arg = builder.CreateLoad(arg_ptr); - args.push_back(arg); - } - builder.CreateCall(var_func, args); - builder.CreateRetVoid(); -} - LogicalResult runLLVMIRPasses(const LLVMTargetOptions& options, llvm::TargetMachine* machine, llvm::Module* module) { @@ -91,7 +53,8 @@ llvm::ModuleAnalysisManager moduleAnalysisManager; llvm::PassInstrumentationCallbacks passInstrumentationCallbacks; - llvm::StandardInstrumentations standardInstrumentations; + llvm::StandardInstrumentations standardInstrumentations( + /*DebugLogging=*/false); standardInstrumentations.registerCallbacks(passInstrumentationCallbacks); llvm::PassBuilder passBuilder(machine, options.pipelineTuningOptions, {},
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/LLVMIRPasses.h b/iree/compiler/Dialect/HAL/Target/LLVM/LLVMIRPasses.h index 199e36f..37ee1ba 100644 --- a/iree/compiler/Dialect/HAL/Target/LLVM/LLVMIRPasses.h +++ b/iree/compiler/Dialect/HAL/Target/LLVM/LLVMIRPasses.h
@@ -31,9 +31,6 @@ std::unique_ptr<llvm::TargetMachine> createTargetMachine( const LLVMTargetOptions& options); -// Creates an invocation function in a module for the given function name. -void createLLVMInvocationFunc(const std::string& name, llvm::Module* module); - // Creates and runs LLVMIR optimization passes defined in LLVMTargetOptions. LogicalResult runLLVMIRPasses(const LLVMTargetOptions& options, llvm::TargetMachine* machine,
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/LLVMIRTarget.cpp b/iree/compiler/Dialect/HAL/Target/LLVM/LLVMIRTarget.cpp index 98c0bf4..96bb5ac 100644 --- a/iree/compiler/Dialect/HAL/Target/LLVM/LLVMIRTarget.cpp +++ b/iree/compiler/Dialect/HAL/Target/LLVM/LLVMIRTarget.cpp
@@ -58,13 +58,9 @@ auto executableOp = cast<IREE::HAL::ExecutableOp>(targetOp.getParentOp()); auto entryPointOps = executableOp.getBlock().getOps<IREE::HAL::ExecutableEntryPointOp>(); - const bool addCInterface = true; for (auto entryPointOp : entryPointOps) { - std::string funcName = - addCInterface ? "_mlir_ciface_" + std::string(entryPointOp.sym_name()) - : std::string(entryPointOp.sym_name()); - llvmIrExecutableDef.entry_points.push_back(funcName); - createLLVMInvocationFunc(funcName, llvmModule.get()); + llvmIrExecutableDef.entry_points.push_back( + std::string(entryPointOp.sym_name())); } // LLVMIR opt passes. @@ -74,8 +70,9 @@ options_.targetTriple); return failure(); } - if (failed( - runLLVMIRPasses(options_, targetMachine.get(), llvmModule.get()))) { + LogicalResult translationResult = + runLLVMIRPasses(options_, targetMachine.get(), llvmModule.get()); + if (failed(translationResult)) { return targetOp.emitError( "Can't build LLVMIR opt passes for ExecutableOp module"); }
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/internal/BUILD b/iree/compiler/Dialect/HAL/Target/LLVM/internal/BUILD index c3ba845..19c3372 100644 --- a/iree/compiler/Dialect/HAL/Target/LLVM/internal/BUILD +++ b/iree/compiler/Dialect/HAL/Target/LLVM/internal/BUILD
@@ -24,5 +24,6 @@ deps = [ "//iree/base:status", "//iree/compiler/Dialect/HAL/Target/LLVM:LLVMAOTTargetLinker_hdrs", + "@llvm-project//llvm:Support", ], )
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/internal/CMakeLists.txt b/iree/compiler/Dialect/HAL/Target/LLVM/internal/CMakeLists.txt index 18c9c7a..b91dae5 100644 --- a/iree/compiler/Dialect/HAL/Target/LLVM/internal/CMakeLists.txt +++ b/iree/compiler/Dialect/HAL/Target/LLVM/internal/CMakeLists.txt
@@ -20,6 +20,7 @@ SRCS "LLVMAOTTargetLinker.cpp" DEPS + LLVMSupport iree::base::status iree::compiler::Dialect::HAL::Target::LLVM::LLVMAOTTargetLinker_hdrs PUBLIC
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/internal/LLVMAOTTargetLinker.cpp b/iree/compiler/Dialect/HAL/Target/LLVM/internal/LLVMAOTTargetLinker.cpp index e5108d6..65d8e22 100644 --- a/iree/compiler/Dialect/HAL/Target/LLVM/internal/LLVMAOTTargetLinker.cpp +++ b/iree/compiler/Dialect/HAL/Target/LLVM/internal/LLVMAOTTargetLinker.cpp
@@ -15,6 +15,7 @@ #include "iree/compiler/Dialect/HAL/Target/LLVM/LLVMAOTTargetLinker.h" #include "iree/base/status.h" +#include "llvm/Support/ToolOutputFile.h" namespace mlir { namespace iree_compiler { @@ -23,18 +24,47 @@ iree::StatusOr<std::string> linkLLVMAOTObjects( const std::string& linkerToolPath, const std::string& objData) { - std::string archiveFile, sharedLibFile; - ASSIGN_OR_RETURN(archiveFile, iree::file_io::GetTempFile("objfile")); - RETURN_IF_ERROR(iree::file_io::SetFileContents(archiveFile, objData)); - ASSIGN_OR_RETURN(sharedLibFile, iree::file_io::GetTempFile("dylibfile")); - std::string linkingCmd = - linkerToolPath + " -shared " + archiveFile + " -o " + sharedLibFile; + llvm::SmallString<32> objFilePath, dylibFilePath; + if (std::error_code error = llvm::sys::fs::createTemporaryFile( + "llvmaot_dylibs", "objfile", objFilePath)) { + return iree::InternalErrorBuilder(IREE_LOC) + << "Failed to generate temporary file for objfile : '" + << error.message() << "'"; + } + if (std::error_code error = llvm::sys::fs::createTemporaryFile( + "llvmaot_dylibs", "dylibfile", dylibFilePath)) { + return iree::InternalErrorBuilder(IREE_LOC) + << "Failed to generate temporary file for dylib : '" + << error.message() << "'"; + } + std::error_code error; + auto outputFile = std::make_unique<llvm::ToolOutputFile>( + objFilePath, error, llvm::sys::fs::F_None); + if (error) { + return iree::InternalErrorBuilder(IREE_LOC) + << "Failed to open temporary objfile '" << objFilePath.c_str() + << "' for dylib : '" << error.message() << "'"; + } + + outputFile->os() << objData; + outputFile->os().flush(); + + auto linkingCmd = + (linkerToolPath + " -shared " + objFilePath + " -o " + dylibFilePath) + .str(); int systemRet = system(linkingCmd.c_str()); if (systemRet != 0) { return iree::InternalErrorBuilder(IREE_LOC) << linkingCmd << " failed with exit code " << systemRet; } - return iree::file_io::GetFileContents(sharedLibFile); + + auto dylibData = llvm::MemoryBuffer::getFile(dylibFilePath); + if (!dylibData) { + return iree::InternalErrorBuilder(IREE_LOC) + << "Failed to read temporary dylib file '" << dylibFilePath.c_str() + << "'"; + } + return dylibData.get()->getBuffer().str(); } iree::StatusOr<std::string> linkLLVMAOTObjectsWithLLDElf(
diff --git a/iree/compiler/Dialect/HAL/hal.imports.mlir b/iree/compiler/Dialect/HAL/hal.imports.mlir index 224bb7b..f3d8c62 100644 --- a/iree/compiler/Dialect/HAL/hal.imports.mlir +++ b/iree/compiler/Dialect/HAL/hal.imports.mlir
@@ -214,6 +214,11 @@ ) -> (i32, i32, i32, i32) attributes {nosideeffects} +// Prints out the content of buffers. +vm.import @buffer_view.trace( + %operands : !vm.ref<!hal.buffer_view> ... +) + //===----------------------------------------------------------------------===// // iree::hal::CommandBuffer //===----------------------------------------------------------------------===//
diff --git a/iree/compiler/Dialect/IREE/Transforms/test/drop_compiler_hints.mlir b/iree/compiler/Dialect/IREE/Transforms/test/drop_compiler_hints.mlir index 9486f96..39c18ce 100644 --- a/iree/compiler/Dialect/IREE/Transforms/test/drop_compiler_hints.mlir +++ b/iree/compiler/Dialect/IREE/Transforms/test/drop_compiler_hints.mlir
@@ -1,6 +1,6 @@ // RUN: iree-opt -split-input-file -iree-drop-compiler-hints %s | IreeFileCheck --implicit-check-not="iree.do_not_optimize" %s -// This file is used as an example in docs/developer_overview.md. +// This file is used as an example in docs/developing_iree/developer_overview.md. // If you move or delete it, please update the documentation accordingly. // CHECK-LABEL: @constant
diff --git a/iree/compiler/Dialect/Shape/IR/ShapeInterface.h b/iree/compiler/Dialect/Shape/IR/ShapeInterface.h index a200bfd..b36890e 100644 --- a/iree/compiler/Dialect/Shape/IR/ShapeInterface.h +++ b/iree/compiler/Dialect/Shape/IR/ShapeInterface.h
@@ -58,7 +58,10 @@ } template <typename BuilderTy, typename... ConstructorArgs> - BuilderTy &make(ConstructorArgs &&... args) { + // TODO(suderman): Re-enable clang-format when new version migrates. + // clang-format off + BuilderTy &make(ConstructorArgs &&...args) { + // clang-format on auto instance = std::make_unique<BuilderTy>(std::forward<ConstructorArgs>(args)...); BuilderTy *unowned = instance.get();
diff --git a/iree/compiler/Dialect/Shape/Plugins/XLA/XlaHloShapeBuilder.cpp b/iree/compiler/Dialect/Shape/Plugins/XLA/XlaHloShapeBuilder.cpp index 8d37751..9024e88 100644 --- a/iree/compiler/Dialect/Shape/Plugins/XLA/XlaHloShapeBuilder.cpp +++ b/iree/compiler/Dialect/Shape/Plugins/XLA/XlaHloShapeBuilder.cpp
@@ -308,6 +308,130 @@ return builder.create<MakeRankedShapeOp>(loc, resultShape, dynamicDims); } +Value rewriteTorchIndexSelect(RankedShapeType resultShape, + TorchIndexSelectOp torchIndexSelectOp, + OpBuilder &builder) { + if (!torchIndexSelectOp) return nullptr; + auto loc = torchIndexSelectOp.getLoc(); + + int64_t resultShapeRank = resultShape.getRank(); + auto paramsType = + torchIndexSelectOp.input().getType().dyn_cast<RankedTensorType>(); + auto indicesType = + torchIndexSelectOp.index().getType().dyn_cast<RankedTensorType>(); + if (!paramsType || !indicesType) { + return nullptr; + } + + auto axis = torchIndexSelectOp.dim(); + auto batchDim = torchIndexSelectOp.batch_dims(); + int64_t paramsRank = paramsType.getRank(); + int64_t indicesRank = indicesType.getRank(); + + std::vector<int64_t> shape(paramsType.getShape()); + int64_t axisValue = axis.getSExtValue(); + int64_t batchDimValue = batchDim.getSExtValue(); + + // For neg axis values, we wrap around params, + // e.g. axis = -1 => params[:-1] + if (axisValue < 0) { + axisValue += paramsRank; + } + if (batchDimValue < 0) { + batchDimValue += indicesRank; + } + + // params must be at least rank axis + 1 + if (paramsRank < axisValue + 1) { + return nullptr; + } + + auto paramsShapeValue = builder.create<GetRankedShapeOp>( + loc, RankedShapeType::get(paramsType.getShape(), builder.getContext()), + torchIndexSelectOp.input()); + auto indicesShapeValue = builder.create<GetRankedShapeOp>( + loc, RankedShapeType::get(indicesType.getShape(), builder.getContext()), + torchIndexSelectOp.index()); + + SmallVector<Value, 4> dynamicDims; +#define GENERATE_RANKED_DIM_OP(value, index) \ + do { \ + auto dimValue = builder.create<RankedDimOp>( \ + loc, builder.getIndexType(), value, builder.getI64IntegerAttr(index)); \ + dynamicDims.push_back(dimValue); \ + } while (0) + + if (indicesRank == 0) { + // Scalar indices (output is rank(params) - 1). + if (resultShapeRank != paramsRank - 1) { + return nullptr; + } + + // params.shape[:axis] + params.shape[axis+1:] + for (int64_t i = 0; i < paramsRank; ++i) { + if ((i == axisValue) || (i < axisValue && !resultShape.isDimDynamic(i)) || + (i > axisValue && !resultShape.isDimDynamic(i - 1))) + continue; + GENERATE_RANKED_DIM_OP(paramsShapeValue, i); + } + } else if (indicesRank == 1) { + // Vector indices (output is rank(params)). + // Copy indices.shape into params.shape[axis] + if (resultShapeRank != paramsRank) { + return nullptr; + } + + // params.shape[:axis] + indices.shape[batch_dims:] + // + params.shape[indicesRank-batchDim+axisValue:] + int resultShapeIndex = 0; + // params.shape[:axis] + for (int64_t i = 0; i < axisValue; ++i) { + if (!resultShape.isDimDynamic(resultShapeIndex++)) continue; + GENERATE_RANKED_DIM_OP(paramsShapeValue, i); + } + // indices.shape[:batchDim] + for (int64_t i = batchDimValue; + i < indicesRank && resultShapeIndex < resultShapeRank; ++i) { + if (!resultShape.isDimDynamic(resultShapeIndex++)) continue; + GENERATE_RANKED_DIM_OP(indicesShapeValue, i); + } + // params.shape[indicesRank-batchDim+axisValue:] + // resultShapeIndex == indicesRank-batchDim+axisValue + for (int64_t i = resultShapeIndex; i < resultShapeRank; ++i) { + if (!resultShape.isDimDynamic(resultShapeIndex++)) continue; + GENERATE_RANKED_DIM_OP(paramsShapeValue, i); + } + } else { + // params.shape[:axis] + indices.shape[batch_dims:] + params.shape[axis + + // 1:] + // The expected rank is (paramsRank-1) + (indicesRank-batchDim) + auto expectedRank = paramsRank - 1 + indicesRank - batchDimValue; + if (resultShapeRank != expectedRank) { + return nullptr; + } + + int resultShapeIndex = 0; + for (int64_t i = 0; i < axisValue; ++i) { + if (!resultShape.isDimDynamic(resultShapeIndex++)) continue; + GENERATE_RANKED_DIM_OP(paramsShapeValue, i); + } + + for (int64_t i = batchDimValue; i < indicesRank; ++i) { + if (!resultShape.isDimDynamic(resultShapeIndex++)) continue; + GENERATE_RANKED_DIM_OP(indicesShapeValue, i); + } + + for (int64_t i = axisValue + 1; + i < paramsRank && resultShapeIndex < resultShapeRank; ++i) { + if (!resultShape.isDimDynamic(resultShapeIndex++)) continue; + GENERATE_RANKED_DIM_OP(paramsShapeValue, i); + } + } +#undef GENERATE_RANKED_DIM_OP + + return builder.create<MakeRankedShapeOp>(loc, resultShape, dynamicDims); +} + } // namespace // Creates a custom op shape builder for XLA-HLO ops that are not otherwise @@ -340,6 +464,8 @@ b.insertOpRankedShapeBuilder<TransposeOp>(rewriteTranspose); b.insertOpRankedShapeBuilder<mhlo::DotGeneralOp>(rewriteDotGeneral); b.insertOpRankedShapeBuilder<mhlo::DynamicReshapeOp>(rewriteDynamicReshape); + b.insertOpRankedShapeBuilder<mhlo::TorchIndexSelectOp>( + rewriteTorchIndexSelect); } } // namespace mhlo
diff --git a/iree/compiler/Dialect/VM/IR/VMBase.td b/iree/compiler/Dialect/VM/IR/VMBase.td index b6c31df..9243cef 100644 --- a/iree/compiler/Dialect/VM/IR/VMBase.td +++ b/iree/compiler/Dialect/VM/IR/VMBase.td
@@ -56,15 +56,39 @@ // VM opcodes //===----------------------------------------------------------------------===// // Opcode ranges: -// 0x00-0x7F: core VM opcodes, reserved for this dialect -// 0x80-0xFF: unreserved, used by target-specific ops (like SIMD) +// 0x00-0x9F: core VM opcodes, reserved for this dialect +// 0xA0-0xFF: unreserved, used to prefix extension op sets // // Note that changing existing opcode assignments will invalidate all binaries // and should only be done when breaking changes are acceptable. We could add a // versioning system here to automatically switch between different encodings // but we are a long way out to stabilizing this format :) +// +// Some opcodes require an extension prefix to indicate that runtime support +// is optional. An op with the ExtI64 trait will require VM_OPC_ExtI64, for +// example. Ops that bridge extension sets have a canonical form that may +// require multiple prefix codes (for example, the i64<->f64 extensions). -class VM_OPC<int opcode, string name> : I32EnumAttrCase<name, opcode>; +class VM_OPC<int opcode, string name> : + IntEnumAttrCaseBase<I8, name, name, opcode>; + +class VM_OPC_EnumAttr<string name, string enumName, string enumTag, + string description, + VM_OPC prefix = ?, + list<VM_OPC> cases> : + IntEnumAttr<I8, name, description, cases> { + let cppNamespace = "IREE::VM"; + let returnType = cppNamespace # "::" # name; + let underlyingType = "uint8_t"; + let convertFromStorage = "static_cast<" # returnType # ">($_self.getInt())"; + let constBuilderCall = + "$_builder.getI8IntegerAttr(static_cast<int8_t>($0))"; + + // Used by VMOpTableGen: + string opcodeEnumName = enumName; + VM_OPC opcodePrefix = prefix; + string opcodeEnumTag = enumTag; +} // Globals: def VM_OPC_GlobalLoadI32 : VM_OPC<0x00, "GlobalLoadI32">; @@ -123,10 +147,12 @@ def VM_OPC_ShrI32U : VM_OPC<0x2F, "ShrI32U">; // Casting and type conversion/emulation: -def VM_OPC_TruncI8 : VM_OPC<0x31, "TruncI8">; -def VM_OPC_TruncI16 : VM_OPC<0x32, "TruncI16">; +def VM_OPC_TruncI32I8 : VM_OPC<0x31, "TruncI32I8">; +def VM_OPC_TruncI32I16 : VM_OPC<0x32, "TruncI32I16">; def VM_OPC_ExtI8I32S : VM_OPC<0x33, "ExtI8I32S">; -def VM_OPC_ExtI16I32S : VM_OPC<0x34, "ExtI16I32S">; +def VM_OPC_ExtI8I32U : VM_OPC<0x34, "ExtI8I32U">; +def VM_OPC_ExtI16I32S : VM_OPC<0x35, "ExtI16I32S">; +def VM_OPC_ExtI16I32U : VM_OPC<0x36, "ExtI16I32U">; // Reduction arithmetic: @@ -135,12 +161,6 @@ def VM_OPC_CmpNEI32 : VM_OPC<0x41, "CmpNEI32">; def VM_OPC_CmpLTI32S : VM_OPC<0x42, "CmpLTI32S">; def VM_OPC_CmpLTI32U : VM_OPC<0x43, "CmpLTI32U">; -def VM_OPC_CmpLTEI32S : VM_OPC<0x44, "CmpLTEI32S">; -def VM_OPC_CmpLTEI32U : VM_OPC<0x45, "CmpLTEI32U">; -def VM_OPC_CmpGTI32S : VM_OPC<0x46, "CmpGTI32S">; -def VM_OPC_CmpGTI32U : VM_OPC<0x47, "CmpGTI32U">; -def VM_OPC_CmpGTEI32S : VM_OPC<0x48, "CmpGTEI32S">; -def VM_OPC_CmpGTEI32U : VM_OPC<0x49, "CmpGTEI32U">; def VM_OPC_CmpNZI32 : VM_OPC<0x4D, "CmpNZI32">; def VM_OPC_CmpEQRef : VM_OPC<0x4A, "CmpEQRef">; def VM_OPC_CmpNERef : VM_OPC<0x4B, "CmpNERef">; @@ -163,8 +183,19 @@ def VM_OPC_CondBreak : VM_OPC<0x7E, "CondBreak">; def VM_OPC_Break : VM_OPC<0x7F, "Break">; -def VM_OpcodeAttr : I32EnumAttr<"Opcode", "valid VM operation encodings", [ - // Core VM opcodes (0x00-0x7F): +// Extension prefixes: +def VM_OPC_PrefixExtI64 : VM_OPC<0xA0, "PrefixExtI64">; +def VM_OPC_PrefixExtF32 : VM_OPC<0xA1, "PrefixExtF32">; +def VM_OPC_PrefixExtF64 : VM_OPC<0xA2, "PrefixExtF64">; + +// Runtime enum iree_vm_core_op_t: +def VM_CoreOpcodeAttr : + VM_OPC_EnumAttr<"Opcode", + "iree_vm_core_op_t", + "CORE", // IREE_VM_OP_CORE_* + "valid VM core operation encodings", + ?, [ + // Core VM opcodes (0x00-0x9F): VM_OPC_GlobalLoadI32, VM_OPC_GlobalStoreI32, VM_OPC_GlobalLoadIndirectI32, @@ -203,16 +234,16 @@ VM_OPC_ShlI32, VM_OPC_ShrI32S, VM_OPC_ShrI32U, + VM_OPC_TruncI32I8, + VM_OPC_TruncI32I16, + VM_OPC_ExtI8I32S, + VM_OPC_ExtI8I32U, + VM_OPC_ExtI16I32S, + VM_OPC_ExtI16I32U, VM_OPC_CmpEQI32, VM_OPC_CmpNEI32, VM_OPC_CmpLTI32S, VM_OPC_CmpLTI32U, - VM_OPC_CmpLTEI32S, - VM_OPC_CmpLTEI32U, - VM_OPC_CmpGTI32S, - VM_OPC_CmpGTI32U, - VM_OPC_CmpGTEI32S, - VM_OPC_CmpGTEI32U, VM_OPC_CmpNZI32, VM_OPC_CmpEQRef, VM_OPC_CmpNERef, @@ -229,11 +260,99 @@ VM_OPC_CondBreak, VM_OPC_Break, - // Extension opcodes (0x80-0xFF): - // TODO(benvanik): SIMD dialect. - ]> { - let cppNamespace = "IREE::VM"; -} + // Extension opcodes (0xA0-0xFF): + VM_OPC_PrefixExtI64, // VM_ExtI64OpcodeAttr + VM_OPC_PrefixExtF32, // VM_ExtF32OpcodeAttr + VM_OPC_PrefixExtF64, // VM_ExtF64OpcodeAttr + ]>; + +// i64 extension: +// (ops are encoded as a VM_OPC_ExtI64 + the opcode below) +def VM_OPC_GlobalLoadI64 : VM_OPC<0x00, "GlobalLoadI64">; +def VM_OPC_GlobalStoreI64 : VM_OPC<0x01, "GlobalStoreI64">; +def VM_OPC_GlobalLoadIndirectI64 : VM_OPC<0x02, "GlobalLoadIndirectI64">; +def VM_OPC_GlobalStoreIndirectI64: VM_OPC<0x03, "GlobalStoreIndirectI64">; +def VM_OPC_ConstI64Zero : VM_OPC<0x08, "ConstI64Zero">; +def VM_OPC_ConstI64 : VM_OPC<0x09, "ConstI64">; +def VM_OPC_ListGetI64 : VM_OPC<0x14, "ListGetI64">; +def VM_OPC_ListSetI64 : VM_OPC<0x15, "ListSetI64">; +def VM_OPC_SelectI64 : VM_OPC<0x1E, "SelectI64">; +def VM_OPC_SwitchI64 : VM_OPC<0x20, "SwitchI64">; +def VM_OPC_AddI64 : VM_OPC<0x22, "AddI64">; +def VM_OPC_SubI64 : VM_OPC<0x23, "SubI64">; +def VM_OPC_MulI64 : VM_OPC<0x24, "MulI64">; +def VM_OPC_DivI64S : VM_OPC<0x25, "DivI64S">; +def VM_OPC_DivI64U : VM_OPC<0x26, "DivI64U">; +def VM_OPC_RemI64S : VM_OPC<0x27, "RemI64S">; +def VM_OPC_RemI64U : VM_OPC<0x28, "RemI64U">; +def VM_OPC_NotI64 : VM_OPC<0x29, "NotI64">; +def VM_OPC_AndI64 : VM_OPC<0x2A, "AndI64">; +def VM_OPC_OrI64 : VM_OPC<0x2B, "OrI64">; +def VM_OPC_XorI64 : VM_OPC<0x2C, "XorI64">; +def VM_OPC_ShlI64 : VM_OPC<0x2D, "ShlI64">; +def VM_OPC_ShrI64S : VM_OPC<0x2E, "ShrI64S">; +def VM_OPC_ShrI64U : VM_OPC<0x2F, "ShrI64U">; +def VM_OPC_TruncI64I8 : VM_OPC<0x30, "TruncI64I8">; +def VM_OPC_TruncI64I16 : VM_OPC<0x31, "TruncI64I16">; +def VM_OPC_TruncI64I32 : VM_OPC<0x32, "TruncI64I32">; +def VM_OPC_ExtI8I64S : VM_OPC<0x33, "ExtI8I64S">; +def VM_OPC_ExtI8I64U : VM_OPC<0x34, "ExtI8I64U">; +def VM_OPC_ExtI16I64S : VM_OPC<0x35, "ExtI16I64S">; +def VM_OPC_ExtI16I64U : VM_OPC<0x36, "ExtI16I64U">; +def VM_OPC_ExtI32I64S : VM_OPC<0x37, "ExtI32I64S">; +def VM_OPC_ExtI32I64U : VM_OPC<0x38, "ExtI32I64U">; +def VM_OPC_CmpEQI64 : VM_OPC<0x40, "CmpEQI64">; +def VM_OPC_CmpNEI64 : VM_OPC<0x41, "CmpNEI64">; +def VM_OPC_CmpLTI64S : VM_OPC<0x42, "CmpLTI64S">; +def VM_OPC_CmpLTI64U : VM_OPC<0x43, "CmpLTI64U">; +def VM_OPC_CmpNZI64 : VM_OPC<0x4D, "CmpNZI64">; + +// Runtime enum iree_vm_ext_i64_op_t: +def VM_ExtI64OpcodeAttr : + VM_OPC_EnumAttr<"ExtI64Opcode", + "iree_vm_ext_i64_op_t", + "EXT_I64", // IREE_VM_OP_EXT_I64_* + "valid VM operation encodings in the i64 extension", + VM_OPC_PrefixExtI64, [ + VM_OPC_GlobalLoadI64, + VM_OPC_GlobalStoreI64, + VM_OPC_GlobalLoadIndirectI64, + VM_OPC_GlobalStoreIndirectI64, + VM_OPC_ConstI64Zero, + VM_OPC_ConstI64, + VM_OPC_ListGetI64, + VM_OPC_ListSetI64, + VM_OPC_SelectI64, + VM_OPC_SwitchI64, + VM_OPC_AddI64, + VM_OPC_SubI64, + VM_OPC_MulI64, + VM_OPC_DivI64S, + VM_OPC_DivI64U, + VM_OPC_RemI64S, + VM_OPC_RemI64U, + VM_OPC_NotI64, + VM_OPC_AndI64, + VM_OPC_OrI64, + VM_OPC_XorI64, + VM_OPC_ShlI64, + VM_OPC_ShrI64S, + VM_OPC_ShrI64U, + VM_OPC_TruncI64I8, + VM_OPC_TruncI64I16, + VM_OPC_TruncI64I32, + VM_OPC_ExtI8I64S, + VM_OPC_ExtI8I64U, + VM_OPC_ExtI16I64S, + VM_OPC_ExtI16I64U, + VM_OPC_ExtI32I64S, + VM_OPC_ExtI32I64U, + VM_OPC_CmpEQI64, + VM_OPC_CmpNEI64, + VM_OPC_CmpLTI64S, + VM_OPC_CmpLTI64U, + VM_OPC_CmpNZI64, + ]>; //===----------------------------------------------------------------------===// // Declarative encoding framework @@ -374,6 +493,40 @@ list<VM_EncEncodeExpr> encoding = ?; } +def VM_GlobalOpInterface : OpInterface<"VMGlobalOp"> { + let description = [{ + Interface used for VM ops that declare global values. + }]; + + let methods = [ + InterfaceMethod<[{ + Returns the storage type of this global such as i32. + }], + "Type", "getStorageType", (ins)>, + InterfaceMethod<[{ + Returns the size in bytes of the global when stored in rwdata. Valid only + for globals using primitive storage. + }], + "size_t", "getStorageSize", (ins), [{ + auto storageType = $_self.getStorageType(); + assert(storageType.isIntOrFloat()); + assert(storageType.getIntOrFloatBitWidth() % 8 == 0); + return storageType.getIntOrFloatBitWidth() / 8; + }]>, + InterfaceMethod<[{}], "StringRef", "getSymbolName", (ins)>, + InterfaceMethod<[{}], "bool", "isMutable", (ins)>, + InterfaceMethod<[{}], "Optional<StringRef>", "getInitializerAttr", (ins)>, + InterfaceMethod<[{}], "Optional<Attribute>", "getInitialValueAttr", (ins)>, + InterfaceMethod<[{}], "Optional<IntegerAttr>", "getOrdinalAttr", (ins)>, + InterfaceMethod<[{}], "int", "getOrdinal", (ins), [{ + return $_self.getOrdinalAttr().getValue().template cast<IntegerAttr>().getInt(); + }]>, + InterfaceMethod<[{}], "void", "makeMutable", (ins)>, + InterfaceMethod<[{}], "void", "clearInitializer", (ins)>, + InterfaceMethod<[{}], "void", "clearInitialValue", (ins)>, + ]; +} + //===----------------------------------------------------------------------===// // VM traits //===----------------------------------------------------------------------===// @@ -390,6 +543,13 @@ // that execution order before and after the barrier op remains the same. def VM_FullBarrier : NativeOpTrait<"IREE::VM::FullBarrier">; +// Operations with this trait require the VM i64 extension. +def VM_ExtI64 : NativeOpTrait<"IREE::VM::ExtI64">; +// Operations with this trait require the VM f32 extension. +def VM_ExtF32 : NativeOpTrait<"IREE::VM::ExtF32">; +// Operations with this trait require the VM f64 extension. +def VM_ExtF64 : NativeOpTrait<"IREE::VM::ExtF64">; + //===----------------------------------------------------------------------===// // ref<T> types //===----------------------------------------------------------------------===// @@ -471,13 +631,16 @@ def VM_AnyType : AnyTypeOf<[ I32, + I64, + F32, + F64, VM_CondValue, VM_AnyRef, ]>; def VM_PrimitiveType : AnyTypeOf<[ - AnyIntOfWidths<[8, 16, 32]>, - FloatOfWidths<[16, 32]>, + AnyIntOfWidths<[8, 16, 32, 64]>, + FloatOfWidths<[16, 32, 64]>, ]>; class VM_ConstIntValueAttr<I type> : Attr<
diff --git a/iree/compiler/Dialect/VM/IR/VMOpFolders.cpp b/iree/compiler/Dialect/VM/IR/VMOpFolders.cpp index 7ae855b..b3ceb17 100644 --- a/iree/compiler/Dialect/VM/IR/VMOpFolders.cpp +++ b/iree/compiler/Dialect/VM/IR/VMOpFolders.cpp
@@ -590,12 +590,12 @@ // Casting and type conversion/emulation //===----------------------------------------------------------------------===// -OpFoldResult TruncI8Op::fold(ArrayRef<Attribute> operands) { +OpFoldResult TruncI32I8Op::fold(ArrayRef<Attribute> operands) { return constFoldUnaryOp<IntegerAttr>( operands, [&](APInt a) { return a.trunc(8).zext(32); }); } -OpFoldResult TruncI16Op::fold(ArrayRef<Attribute> operands) { +OpFoldResult TruncI32I16Op::fold(ArrayRef<Attribute> operands) { return constFoldUnaryOp<IntegerAttr>( operands, [&](APInt a) { return a.trunc(16).zext(32); }); } @@ -605,11 +605,21 @@ operands, [&](APInt a) { return a.trunc(8).sext(32); }); } +OpFoldResult ExtI8I32UOp::fold(ArrayRef<Attribute> operands) { + return constFoldUnaryOp<IntegerAttr>( + operands, [&](APInt a) { return a.trunc(8).zext(32); }); +} + OpFoldResult ExtI16I32SOp::fold(ArrayRef<Attribute> operands) { return constFoldUnaryOp<IntegerAttr>( operands, [&](APInt a) { return a.trunc(16).sext(32); }); } +OpFoldResult ExtI16I32UOp::fold(ArrayRef<Attribute> operands) { + return constFoldUnaryOp<IntegerAttr>( + operands, [&](APInt a) { return a.trunc(16).zext(32); }); +} + //===----------------------------------------------------------------------===// // Native reduction (horizontal) arithmetic //===----------------------------------------------------------------------===// @@ -706,9 +716,7 @@ } void CmpLTI32SOp::getCanonicalizationPatterns(OwningRewritePatternList &results, - MLIRContext *context) { - results.insert<SwapInvertedCmpOps<CmpLTI32SOp, CmpGTEI32SOp>>(context); -} + MLIRContext *context) {} OpFoldResult CmpLTI32UOp::fold(ArrayRef<Attribute> operands) { if (lhs() == rhs()) { @@ -720,9 +728,27 @@ } void CmpLTI32UOp::getCanonicalizationPatterns(OwningRewritePatternList &results, - MLIRContext *context) { - results.insert<SwapInvertedCmpOps<CmpLTI32UOp, CmpGTEI32UOp>>(context); -} + MLIRContext *context) {} + +namespace { + +/// Rewrites a vm.cmp.lte.* pseudo op to a vm.cmp.lt.* op. +template <typename T, typename U> +struct RewritePseudoCmpLTEToLT : public OpRewritePattern<T> { + using OpRewritePattern<T>::OpRewritePattern; + LogicalResult matchAndRewrite(T op, + PatternRewriter &rewriter) const override { + // !(lhs > rhs) + auto condValue = + rewriter.createOrFold<U>(op.getLoc(), op.getType(), op.rhs(), op.lhs()); + rewriter.replaceOpWithNewOp<XorI32Op>( + op, op.getType(), condValue, + rewriter.createOrFold<IREE::VM::ConstI32Op>(op.getLoc(), 1)); + return success(); + } +}; + +} // namespace OpFoldResult CmpLTEI32SOp::fold(ArrayRef<Attribute> operands) { if (lhs() == rhs()) { @@ -736,6 +762,7 @@ void CmpLTEI32SOp::getCanonicalizationPatterns( OwningRewritePatternList &results, MLIRContext *context) { results.insert<SwapInvertedCmpOps<CmpLTEI32SOp, CmpGTI32SOp>>(context); + results.insert<RewritePseudoCmpLTEToLT<CmpLTEI32SOp, CmpLTI32SOp>>(context); } OpFoldResult CmpLTEI32UOp::fold(ArrayRef<Attribute> operands) { @@ -750,8 +777,25 @@ void CmpLTEI32UOp::getCanonicalizationPatterns( OwningRewritePatternList &results, MLIRContext *context) { results.insert<SwapInvertedCmpOps<CmpLTEI32UOp, CmpGTI32UOp>>(context); + results.insert<RewritePseudoCmpLTEToLT<CmpLTEI32UOp, CmpLTI32UOp>>(context); } +namespace { + +/// Rewrites a vm.cmp.gt.* pseudo op to a vm.cmp.lt.* op. +template <typename T, typename U> +struct RewritePseudoCmpGTToLT : public OpRewritePattern<T> { + using OpRewritePattern<T>::OpRewritePattern; + LogicalResult matchAndRewrite(T op, + PatternRewriter &rewriter) const override { + // rhs < lhs + rewriter.replaceOpWithNewOp<U>(op, op.getType(), op.rhs(), op.lhs()); + return success(); + } +}; + +} // namespace + OpFoldResult CmpGTI32SOp::fold(ArrayRef<Attribute> operands) { if (lhs() == rhs()) { // x > x = false @@ -764,6 +808,7 @@ void CmpGTI32SOp::getCanonicalizationPatterns(OwningRewritePatternList &results, MLIRContext *context) { results.insert<SwapInvertedCmpOps<CmpGTI32SOp, CmpLTEI32SOp>>(context); + results.insert<RewritePseudoCmpGTToLT<CmpGTI32SOp, CmpLTI32SOp>>(context); } OpFoldResult CmpGTI32UOp::fold(ArrayRef<Attribute> operands) { @@ -778,8 +823,29 @@ void CmpGTI32UOp::getCanonicalizationPatterns(OwningRewritePatternList &results, MLIRContext *context) { results.insert<SwapInvertedCmpOps<CmpGTI32UOp, CmpLTEI32UOp>>(context); + results.insert<RewritePseudoCmpGTToLT<CmpGTI32UOp, CmpLTI32UOp>>(context); } +namespace { + +/// Rewrites a vm.cmp.gte.* pseudo op to a vm.cmp.lt.* op. +template <typename T, typename U> +struct RewritePseudoCmpGTEToLT : public OpRewritePattern<T> { + using OpRewritePattern<T>::OpRewritePattern; + LogicalResult matchAndRewrite(T op, + PatternRewriter &rewriter) const override { + // !(lhs < rhs) + auto condValue = + rewriter.createOrFold<U>(op.getLoc(), op.getType(), op.lhs(), op.rhs()); + rewriter.replaceOpWithNewOp<XorI32Op>( + op, op.getType(), condValue, + rewriter.createOrFold<IREE::VM::ConstI32Op>(op.getLoc(), 1)); + return success(); + } +}; + +} // namespace + OpFoldResult CmpGTEI32SOp::fold(ArrayRef<Attribute> operands) { if (lhs() == rhs()) { // x >= x = true @@ -792,6 +858,7 @@ void CmpGTEI32SOp::getCanonicalizationPatterns( OwningRewritePatternList &results, MLIRContext *context) { results.insert<SwapInvertedCmpOps<CmpGTEI32SOp, CmpLTI32SOp>>(context); + results.insert<RewritePseudoCmpGTEToLT<CmpGTEI32SOp, CmpLTI32SOp>>(context); } OpFoldResult CmpGTEI32UOp::fold(ArrayRef<Attribute> operands) { @@ -806,6 +873,7 @@ void CmpGTEI32UOp::getCanonicalizationPatterns( OwningRewritePatternList &results, MLIRContext *context) { results.insert<SwapInvertedCmpOps<CmpGTEI32UOp, CmpLTI32UOp>>(context); + results.insert<RewritePseudoCmpGTEToLT<CmpGTEI32UOp, CmpLTI32UOp>>(context); } OpFoldResult CmpNZI32Op::fold(ArrayRef<Attribute> operands) {
diff --git a/iree/compiler/Dialect/VM/IR/VMOps.td b/iree/compiler/Dialect/VM/IR/VMOps.td index c552347..b3a7a4e 100644 --- a/iree/compiler/Dialect/VM/IR/VMOps.td +++ b/iree/compiler/Dialect/VM/IR/VMOps.td
@@ -256,6 +256,7 @@ IsolatedFromAbove, HasParent<"IREE::VM::ModuleOp">, Symbol, + VM_GlobalOpInterface, ])> { let arguments = (ins StrAttr:$sym_name, @@ -291,9 +292,15 @@ ]; let extraClassDeclaration = [{ + StringRef getSymbolName() { return sym_name(); } + Type getStorageType() { return type(); } + bool isMutable() { return is_mutable(); } void makeMutable() { setAttr("is_mutable", UnitAttr::get(getContext())); } + Optional<StringRef> getInitializerAttr() { return initializer(); } void clearInitializer() { removeAttr("initializer"); } + Optional<Attribute> getInitialValueAttr() { return initial_valueAttr(); } void clearInitialValue() { removeAttr("initial_value"); } + Optional<IntegerAttr> getOrdinalAttr() { return ordinalAttr(); } }]; let parser = [{ return parseGlobalOp(parser, &result); }]; @@ -1248,7 +1255,7 @@ let encoding = [ VM_EncOpcode<opcode>, VM_EncOperand<"operand", 0>, - VM_EncIntAttr<"amount", type.bitwidth>, + VM_EncIntAttr<"amount", 8>, VM_EncResult<"result">, ]; } @@ -1272,12 +1279,12 @@ // Casting and type conversion/emulation //===----------------------------------------------------------------------===// -def VM_TruncI8Op : VM_UnaryArithmeticOp<I32, "trunc.i8", VM_OPC_TruncI8> { +def VM_TruncI32I8Op : VM_UnaryArithmeticOp<I32, "trunc.i32.i8", VM_OPC_TruncI32I8> { let summary = [{integer truncate to 8 bits}]; let hasFolder = 1; } -def VM_TruncI16Op : VM_UnaryArithmeticOp<I32, "trunc.i16", VM_OPC_TruncI16> { +def VM_TruncI32I16Op : VM_UnaryArithmeticOp<I32, "trunc.i32.i16", VM_OPC_TruncI32I16> { let summary = [{integer truncate to 16 bits}]; let hasFolder = 1; } @@ -1287,11 +1294,21 @@ let hasFolder = 1; } +def VM_ExtI8I32UOp : VM_UnaryArithmeticOp<I32, "ext.i8.i32.u", VM_OPC_ExtI8I32U> { + let summary = [{integer zero extend 8 bits to 32 bits}]; + let hasFolder = 1; +} + def VM_ExtI16I32SOp : VM_UnaryArithmeticOp<I32, "ext.i16.i32.s", VM_OPC_ExtI16I32S> { let summary = [{integer sign extend 16 bits to 32 bits}]; let hasFolder = 1; } +def VM_ExtI16I32UOp : VM_UnaryArithmeticOp<I32, "ext.i16.i32.u", VM_OPC_ExtI16I32U> { + let summary = [{integer zero extend 16 bits to 32 bits}]; + let hasFolder = 1; +} + //===----------------------------------------------------------------------===// // Native reduction (horizontal) arithmetic //===----------------------------------------------------------------------===// @@ -1355,6 +1372,27 @@ ]; } +class VM_BinaryComparisonPseudoOp<Type type, string mnemonic, + list<OpTrait> traits = []> : + VM_PureOp<mnemonic, !listconcat(traits, [ + AllTypesMatch<["lhs", "rhs"]>, + VM_PseudoOp, + ])> { + let description = [{ + Compares two operands with the specified predicate. + }]; + + let arguments = (ins + type:$lhs, + type:$rhs + ); + let results = (outs + I32:$result + ); + + let assemblyFormat = "operands attr-dict `:` type($lhs)"; +} + def VM_CmpEQI32Op : VM_BinaryComparisonOp<I32, "cmp.eq.i32", VM_OPC_CmpEQI32, [Commutative]> { let summary = [{integer equality comparison operation}]; @@ -1384,43 +1422,42 @@ } def VM_CmpLTEI32SOp : - VM_BinaryComparisonOp<I32, "cmp.lte.i32.s", VM_OPC_CmpLTEI32S> { + VM_BinaryComparisonPseudoOp<I32, "cmp.lte.i32.s"> { let summary = [{signed integer less-than-or-equal comparison operation}]; let hasCanonicalizer = 1; let hasFolder = 1; } def VM_CmpLTEI32UOp : - VM_BinaryComparisonOp<I32, "cmp.lte.i32.u", VM_OPC_CmpLTEI32U> { + VM_BinaryComparisonPseudoOp<I32, "cmp.lte.i32.u"> { let summary = [{unsigned integer less-than-or-equal comparison operation}]; let hasCanonicalizer = 1; let hasFolder = 1; } -// TODO(benvanik): drop these and rely on lt/lte only? def VM_CmpGTI32SOp : - VM_BinaryComparisonOp<I32, "cmp.gt.i32.s", VM_OPC_CmpGTI32S> { + VM_BinaryComparisonPseudoOp<I32, "cmp.gt.i32.s"> { let summary = [{signed integer greater-than comparison operation}]; let hasCanonicalizer = 1; let hasFolder = 1; } def VM_CmpGTI32UOp : - VM_BinaryComparisonOp<I32, "cmp.gt.i32.u", VM_OPC_CmpGTI32U> { + VM_BinaryComparisonPseudoOp<I32, "cmp.gt.i32.u"> { let summary = [{unsigned integer greater-than comparison operation}]; let hasCanonicalizer = 1; let hasFolder = 1; } def VM_CmpGTEI32SOp : - VM_BinaryComparisonOp<I32, "cmp.gte.i32.s", VM_OPC_CmpGTEI32S> { + VM_BinaryComparisonPseudoOp<I32, "cmp.gte.i32.s"> { let summary = [{signed integer greater-than-or-equal comparison operation}]; let hasCanonicalizer = 1; let hasFolder = 1; } def VM_CmpGTEI32UOp : - VM_BinaryComparisonOp<I32, "cmp.gte.i32.u", VM_OPC_CmpGTEI32U> { + VM_BinaryComparisonPseudoOp<I32, "cmp.gte.i32.u"> { let summary = [{unsigned integer greater-than-or-equal comparison operation}]; let hasCanonicalizer = 1; let hasFolder = 1;
diff --git a/iree/compiler/Dialect/VM/IR/VMTraits.h b/iree/compiler/Dialect/VM/IR/VMTraits.h index 7527383..330f75c 100644 --- a/iree/compiler/Dialect/VM/IR/VMTraits.h +++ b/iree/compiler/Dialect/VM/IR/VMTraits.h
@@ -49,6 +49,33 @@ } }; +template <typename ConcreteType> +class ExtI64 : public OpTrait::TraitBase<ConcreteType, ExtI64> { + public: + static LogicalResult verifyTrait(Operation *op) { + // TODO(benvanik): verify i64 ext is supported. + return success(); + } +}; + +template <typename ConcreteType> +class ExtF32 : public OpTrait::TraitBase<ConcreteType, ExtF32> { + public: + static LogicalResult verifyTrait(Operation *op) { + // TODO(benvanik): verify f32 ext is supported. + return success(); + } +}; + +template <typename ConcreteType> +class ExtF64 : public OpTrait::TraitBase<ConcreteType, ExtF64> { + public: + static LogicalResult verifyTrait(Operation *op) { + // TODO(benvanik): verify f64 ext is supported. + return success(); + } +}; + } // namespace VM } // namespace IREE } // namespace OpTrait
diff --git a/iree/compiler/Dialect/VM/IR/test/conversion_folding.mlir b/iree/compiler/Dialect/VM/IR/test/conversion_folding.mlir index 0f6222e..02a9ea3 100644 --- a/iree/compiler/Dialect/VM/IR/test/conversion_folding.mlir +++ b/iree/compiler/Dialect/VM/IR/test/conversion_folding.mlir
@@ -8,7 +8,7 @@ vm.func @trunc_i8_const() -> i32 { // CHECK: vm.const.i32 255 : i32 %c = vm.const.i32 0xFFFFFFFF : i32 - %0 = vm.trunc.i8 %c : i32 + %0 = vm.trunc.i32.i8 %c : i32 vm.return %0 : i32 } @@ -16,7 +16,7 @@ vm.func @trunc_i16_const() -> i32 { // CHECK: vm.const.i32 65535 : i32 %c = vm.const.i32 0xFFFFFFFF : i32 - %0 = vm.trunc.i16 %c : i32 + %0 = vm.trunc.i32.i16 %c : i32 vm.return %0 : i32 } } @@ -33,6 +33,14 @@ vm.return %0 : i32 } + // CHECK-LABEL: @ext_i8_i32_u_const + vm.func @ext_i8_i32_u_const() -> i32 { + // CHECK: vm.const.i32 255 : i32 + %c = vm.const.i32 0x000000FF : i32 + %0 = vm.ext.i8.i32.u %c : i32 + vm.return %0 : i32 + } + // CHECK-LABEL: @ext_i16_i32_s_const vm.func @ext_i16_i32_s_const() -> i32 { // CHECK: vm.const.i32 -1 : i32 @@ -40,4 +48,12 @@ %0 = vm.ext.i16.i32.s %c : i32 vm.return %0 : i32 } + + // CHECK-LABEL: @ext_i16_i32_u_const + vm.func @ext_i16_i32_u_const() -> i32 { + // CHECK: vm.const.i32 65535 : i32 + %c = vm.const.i32 0x0000FFFF : i32 + %0 = vm.ext.i16.i32.u %c : i32 + vm.return %0 : i32 + } }
diff --git a/iree/compiler/Dialect/VM/IR/test/conversion_ops.mlir b/iree/compiler/Dialect/VM/IR/test/conversion_ops.mlir index 5e94458..008c5d6 100644 --- a/iree/compiler/Dialect/VM/IR/test/conversion_ops.mlir +++ b/iree/compiler/Dialect/VM/IR/test/conversion_ops.mlir
@@ -5,10 +5,10 @@ // CHECK-LABEL: @trunc vm.module @my_module { vm.func @trunc(%arg0 : i32) -> i32 { - // CHECK: %0 = vm.trunc.i8 %arg0 : i32 - %0 = vm.trunc.i8 %arg0 : i32 - // CHECK-NEXT: %1 = vm.trunc.i16 %0 : i32 - %1 = vm.trunc.i16 %0 : i32 + // CHECK: %0 = vm.trunc.i32.i8 %arg0 : i32 + %0 = vm.trunc.i32.i8 %arg0 : i32 + // CHECK-NEXT: %1 = vm.trunc.i32.i16 %0 : i32 + %1 = vm.trunc.i32.i16 %0 : i32 vm.return %1 : i32 } } @@ -20,8 +20,12 @@ vm.func @ext(%arg0 : i32) -> i32 { // CHECK-NEXT: %0 = vm.ext.i8.i32.s %arg0 : i32 %0 = vm.ext.i8.i32.s %arg0 : i32 - // CHECK-NEXT: %1 = vm.ext.i16.i32.s %0 : i32 - %1 = vm.ext.i16.i32.s %0 : i32 - vm.return %1 : i32 + // CHECK-NEXT: %1 = vm.ext.i8.i32.u %0 : i32 + %1 = vm.ext.i8.i32.u %0 : i32 + // CHECK-NEXT: %2 = vm.ext.i16.i32.s %1 : i32 + %2 = vm.ext.i16.i32.s %1 : i32 + // CHECK-NEXT: %3 = vm.ext.i16.i32.u %2 : i32 + %3 = vm.ext.i16.i32.u %2 : i32 + vm.return %3 : i32 } }
diff --git a/iree/compiler/Dialect/VM/Target/Bytecode/BytecodeModuleTarget.cpp b/iree/compiler/Dialect/VM/Target/Bytecode/BytecodeModuleTarget.cpp index 314c6ea..2273286 100644 --- a/iree/compiler/Dialect/VM/Target/Bytecode/BytecodeModuleTarget.cpp +++ b/iree/compiler/Dialect/VM/Target/Bytecode/BytecodeModuleTarget.cpp
@@ -56,7 +56,7 @@ int importFuncs = 0; int exportFuncs = 0; int internalFuncs = 0; - int globalBytes = 0; + size_t globalBytes = 0; int globalRefs = 0; int rodatas = 0; int rwdatas = 0; @@ -85,12 +85,14 @@ ++counts.exportFuncs; } else if (isa<IREE::VM::ImportOp>(op)) { ++counts.importFuncs; - } else if (isa<IREE::VM::GlobalI32Op>(op)) { - counts.globalBytes += 4; - } else if (isa<IREE::VM::GlobalRefOp>(op)) { - ++counts.globalRefs; } else if (isa<IREE::VM::RodataOp>(op)) { ++counts.rodatas; + } else if (isa<IREE::VM::GlobalRefOp>(op)) { + ++counts.globalRefs; + } else if (auto globalOp = dyn_cast<VMGlobalOp>(op)) { + counts.globalBytes = + std::max(counts.globalBytes, + globalOp.getOrdinal() + globalOp.getStorageSize()); } } return counts;
diff --git a/iree/compiler/Dialect/VM/Tools/VMOpTableGen.cpp b/iree/compiler/Dialect/VM/Tools/VMOpTableGen.cpp index de17975..b3081c1 100644 --- a/iree/compiler/Dialect/VM/Tools/VMOpTableGen.cpp +++ b/iree/compiler/Dialect/VM/Tools/VMOpTableGen.cpp
@@ -29,47 +29,33 @@ using ::llvm::formatv; using ::llvm::Record; -// Finds all serializable ops and emits a enum and template table for their -// opcode and name. -bool emitOpTableDefs(const llvm::RecordKeeper &recordKeeper, raw_ostream &os) { - llvm::emitSourceFileHeader("IREE VM Operation Tables", os); - - std::vector<const Record *> opRecords(256); +void emitOpTable(const llvm::RecordKeeper &recordKeeper, const Record &tableDef, + raw_ostream &os) { std::vector<const Record *> opEncodings(256); - auto defs = recordKeeper.getAllDerivedDefinitions("VM_Op"); - for (const auto *def : defs) { - if (def->isValueUnset("encoding")) continue; - auto encodingExprs = def->getValueAsListOfDefs("encoding"); - for (auto encodingExpr : encodingExprs) { - if (encodingExpr->getType()->getAsString() == "VM_EncOpcode") { - auto *opcode = encodingExpr->getValueAsDef("opcode"); - opRecords[opcode->getValueAsInt("value")] = def; - opEncodings[opcode->getValueAsInt("value")] = opcode; - break; - } - } + for (auto *opcodeDef : tableDef.getValueAsListOfDefs("enumerants")) { + opEncodings[opcodeDef->getValueAsInt("value")] = opcodeDef; } os << "typedef enum {\n"; for (int i = 0; i < 256; ++i) { - auto *def = opRecords[i]; - if (def) { - auto *opcode = opEncodings[i]; - os << formatv(" IREE_VM_OP_{0} = {1}", + if (auto *opcode = opEncodings[i]) { + os << formatv(" IREE_VM_OP_{0}_{1} = {2}", + tableDef.getValueAsString("opcodeEnumTag"), opcode->getValueAsString("symbol"), format_hex(i, 4, true)); } else { - os << formatv(" IREE_VM_OP_RSV_{0}", format_hex(i, 4, true)); + os << formatv(" IREE_VM_OP_{0}_RSV_{1}", + tableDef.getValueAsString("opcodeEnumTag"), + format_hex(i, 4, true)); } os << ",\n"; } - os << "} iree_vm_op_t;\n"; + os << "} " << tableDef.getValueAsString("opcodeEnumName") << ";\n"; os << "\n"; - os << "#define IREE_VM_OP_TABLE(OPC, RSV) \\\n"; + os << formatv("#define IREE_VM_OP_{0}_TABLE(OPC, RSV) \\\n", + tableDef.getValueAsString("opcodeEnumTag")); for (int i = 0; i < 256; ++i) { - auto *def = opRecords[i]; - if (def) { - auto *opcode = opEncodings[i]; + if (auto *opcode = opEncodings[i]) { os << formatv(" OPC({0}, {1})", format_hex(i, 4, true), opcode->getValueAsString("symbol")); } else { @@ -80,6 +66,17 @@ } } os << "\n\n"; +} + +// Finds all opcode tables in VMBase.td and emits a enum and template table for +// their opcode and name. +bool emitOpTableDefs(const llvm::RecordKeeper &recordKeeper, raw_ostream &os) { + llvm::emitSourceFileHeader("IREE VM Operation Tables", os); + + auto defs = recordKeeper.getAllDerivedDefinitions("VM_OPC_EnumAttr"); + for (const auto *def : defs) { + emitOpTable(recordKeeper, *def, os); + } return false; }
diff --git a/iree/compiler/Dialect/VM/Transforms/GlobalInitialization.cpp b/iree/compiler/Dialect/VM/Transforms/GlobalInitialization.cpp index b87441d..6a30818 100644 --- a/iree/compiler/Dialect/VM/Transforms/GlobalInitialization.cpp +++ b/iree/compiler/Dialect/VM/Transforms/GlobalInitialization.cpp
@@ -65,13 +65,13 @@ // could gather the ops, sort them (by some rule), and then build the // initialization function. for (auto &op : getOperation().getBlock().getOperations()) { - if (auto globalOp = dyn_cast<GlobalI32Op>(op)) { - if (failed(appendInitialization(globalOp, initBuilder))) { + if (auto globalOp = dyn_cast<GlobalRefOp>(op)) { + if (failed(appendRefInitialization(globalOp, initBuilder))) { globalOp.emitOpError() << "unable to be initialized"; return signalPassFailure(); } - } else if (auto globalOp = dyn_cast<GlobalRefOp>(op)) { - if (failed(appendInitialization(globalOp, initBuilder))) { + } else if (auto globalOp = dyn_cast<VMGlobalOp>(op)) { + if (failed(appendPrimitiveInitialization(globalOp, initBuilder))) { globalOp.emitOpError() << "unable to be initialized"; return signalPassFailure(); } @@ -96,27 +96,72 @@ } private: - LogicalResult appendInitialization(GlobalI32Op globalOp, OpBuilder &builder) { - if (globalOp.initial_value().hasValue()) { - auto constOp = builder.create<ConstI32Op>(globalOp.getLoc(), - globalOp.initial_valueAttr()); - builder.create<GlobalStoreI32Op>(globalOp.getLoc(), constOp.getResult(), - globalOp.sym_name()); + LogicalResult appendPrimitiveInitialization(VMGlobalOp globalOp, + OpBuilder &builder) { + auto initialValue = + globalOp.getInitialValueAttr().getValueOr<Attribute>({}); + Value value = {}; + if (initialValue) { + LogicalResult constResult = success(); + std::tie(constResult, value) = + createConst(globalOp.getLoc(), initialValue, builder); + if (failed(constResult)) { + return globalOp.emitOpError() + << "unable to create initializer constant for global"; + } globalOp.clearInitialValue(); - globalOp.makeMutable(); - } else if (globalOp.initializer().hasValue()) { + } else if (globalOp.getInitializerAttr().hasValue()) { auto callOp = builder.create<CallOp>( - globalOp.getLoc(), globalOp.initializerAttr(), - ArrayRef<Type>{globalOp.type()}, ArrayRef<Value>{}); - builder.create<GlobalStoreI32Op>(globalOp.getLoc(), callOp.getResult(0), - globalOp.sym_name()); + globalOp.getLoc(), globalOp.getInitializerAttr().getValue(), + ArrayRef<Type>{globalOp.getStorageType()}, ArrayRef<Value>{}); + value = callOp.getResult(0); globalOp.clearInitializer(); - globalOp.makeMutable(); } - return success(); + if (!value) { + // Globals are zero-initialized by default so we can just strip the + // initial value/initializer and avoid the work entirely. + return success(); + } + globalOp.makeMutable(); + return storePrimitiveGlobal(globalOp.getLoc(), globalOp.getSymbolName(), + value, builder); } - LogicalResult appendInitialization(GlobalRefOp globalOp, OpBuilder &builder) { + // Returns {} if the constant is zero. + std::pair<LogicalResult, Value> createConst(Location loc, Attribute value, + OpBuilder &builder) { + if (auto intValue = value.dyn_cast<IntegerAttr>()) { + if (intValue.getValue().isNullValue()) { + // Globals are zero-initialized by default. + return {success(), {}}; + } + switch (intValue.getValue().getBitWidth()) { + case 32: + return {success(), builder.createOrFold<ConstI32Op>(loc, intValue)}; + default: + return {failure(), {}}; + } + } + return {failure(), {}}; + } + + // Stores a value to a global; the global must be mutable. + LogicalResult storePrimitiveGlobal(Location loc, StringRef symName, + Value value, OpBuilder &builder) { + if (auto intType = value.getType().dyn_cast<IntegerType>()) { + switch (intType.getIntOrFloatBitWidth()) { + case 32: + builder.create<GlobalStoreI32Op>(loc, value, symName); + return success(); + default: + return failure(); + } + } + return failure(); + } + + LogicalResult appendRefInitialization(GlobalRefOp globalOp, + OpBuilder &builder) { if (globalOp.initializer().hasValue()) { auto callOp = builder.create<CallOp>( globalOp.getLoc(), globalOp.initializerAttr(),
diff --git a/iree/compiler/Dialect/VM/Transforms/OrdinalAllocation.cpp b/iree/compiler/Dialect/VM/Transforms/OrdinalAllocation.cpp index e75b918..6692c7a 100644 --- a/iree/compiler/Dialect/VM/Transforms/OrdinalAllocation.cpp +++ b/iree/compiler/Dialect/VM/Transforms/OrdinalAllocation.cpp
@@ -49,9 +49,9 @@ int nextFuncOrdinal = 0; int nextImportOrdinal = 0; int nextExportOrdinal = 0; - int nextGlobalBytesOrdinal = 0; int nextGlobalRefOrdinal = 0; int nextRodataOrdinal = 0; + SmallVector<SmallVector<VMGlobalOp, 4>, 8> primitiveGlobalOps(8); for (auto &op : getOperation().getBlock().getOperations()) { Optional<int> ordinal = llvm::None; if (auto funcOp = dyn_cast<FuncOp>(op)) { @@ -60,19 +60,36 @@ ordinal = nextExportOrdinal++; } else if (isa<ImportOp>(op)) { ordinal = nextImportOrdinal++; - } else if (isa<GlobalI32Op>(op)) { - ordinal = nextGlobalBytesOrdinal; - nextGlobalBytesOrdinal += 4; - } else if (isa<GlobalRefOp>(op)) { - ordinal = nextGlobalRefOrdinal++; } else if (isa<RodataOp>(op)) { ordinal = nextRodataOrdinal++; + } else if (isa<GlobalRefOp>(op)) { + ordinal = nextGlobalRefOrdinal++; + } else if (auto globalOp = dyn_cast<VMGlobalOp>(op)) { + // Bucket the primitive global ops (like vm.global.i32) so we can + // run over all of them below. + primitiveGlobalOps[globalOp.getStorageSize()].push_back(globalOp); + continue; } if (ordinal.hasValue()) { op.setAttr("ordinal", builder.getI32IntegerAttr(ordinal.getValue())); } } + // Assign byte offset values to primitive globals, ensuring that we meet + // natural alignment requirements on each size type. + int nextGlobalBytesOrdinal = 0; + for (auto sizeGlobalOps : llvm::enumerate(primitiveGlobalOps)) { + size_t storageSize = sizeGlobalOps.index(); + if (sizeGlobalOps.value().empty()) continue; + nextGlobalBytesOrdinal = + llvm::alignTo(nextGlobalBytesOrdinal, storageSize); + for (auto &globalOp : sizeGlobalOps.value()) { + globalOp.setAttr("ordinal", + builder.getI32IntegerAttr(nextGlobalBytesOrdinal)); + nextGlobalBytesOrdinal += storageSize; + } + } + SymbolTable symbolTable(getOperation()); // Convert all global address pseudo-ops to constants referencing the
diff --git a/iree/hal/BUILD b/iree/hal/BUILD index f976cf0..ca34717 100644 --- a/iree/hal/BUILD +++ b/iree/hal/BUILD
@@ -157,7 +157,6 @@ "//iree/base:bitfield", "//iree/base:status", "//iree/base:time", - "@com_google_absl//absl/time", "@com_google_absl//absl/types:span", ], ) @@ -231,7 +230,6 @@ "//iree/base:status", "//iree/base:target_platform", "//iree/base:time", - "@com_google_absl//absl/time", ], ) @@ -262,7 +260,6 @@ "//iree/base:time", "//iree/base:tracing", "@com_google_absl//absl/synchronization", - "@com_google_absl//absl/time", "@com_google_absl//absl/types:span", ], ) @@ -376,7 +373,6 @@ ":resource", "//iree/base:status", "//iree/base:time", - "@com_google_absl//absl/time", ], )
diff --git a/iree/hal/CMakeLists.txt b/iree/hal/CMakeLists.txt index 079f76e..5d22e3c 100644 --- a/iree/hal/CMakeLists.txt +++ b/iree/hal/CMakeLists.txt
@@ -163,7 +163,6 @@ ::command_buffer ::semaphore absl::span - absl::time iree::base::bitfield iree::base::status iree::base::time @@ -247,7 +246,6 @@ ::executable_cache ::executable_layout ::semaphore - absl::time iree::base::ref_ptr iree::base::status iree::base::target_platform @@ -284,7 +282,6 @@ ::semaphore absl::span absl::synchronization - absl::time iree::base::source_location iree::base::status iree::base::time @@ -432,7 +429,6 @@ "semaphore.h" DEPS ::resource - absl::time iree::base::status iree::base::time PUBLIC
diff --git a/iree/hal/api.cc b/iree/hal/api.cc index 088f055..c5dd130 100644 --- a/iree/hal/api.cc +++ b/iree/hal/api.cc
@@ -1795,11 +1795,11 @@ switch (wait_mode) { case IREE_HAL_WAIT_MODE_ALL: wait_status = - handle->WaitAllSemaphores(semaphore_values, ToAbslTime(deadline_ns)); + handle->WaitAllSemaphores(semaphore_values, Time(deadline_ns)); break; case IREE_HAL_WAIT_MODE_ANY: wait_status = std::move(handle->WaitAnySemaphore(semaphore_values, - ToAbslTime(deadline_ns))) + Time(deadline_ns))) .status(); break; default: @@ -1819,8 +1819,7 @@ iree_hal_device_t* device, iree_hal_wait_mode_t wait_mode, const iree_hal_semaphore_list_t* semaphore_list, iree_duration_t timeout_ns) { - iree_time_t deadline_ns = - FromAbslTime(iree::RelativeTimeoutToDeadline(ToAbslDuration(timeout_ns))); + iree_time_t deadline_ns = iree_relative_timeout_to_deadline_ns(timeout_ns); return iree_hal_device_wait_semaphores_with_deadline( device, wait_mode, semaphore_list, deadline_ns); } @@ -2161,7 +2160,7 @@ IREE_TRACE_SCOPE0("iree_hal_semaphore_wait_with_deadline"); auto* handle = reinterpret_cast<Semaphore*>(semaphore); if (!handle) return IREE_STATUS_INVALID_ARGUMENT; - return ToApiStatus(handle->Wait(value, ToAbslTime(deadline_ns))); + return ToApiStatus(handle->Wait(value, Time(deadline_ns))); } IREE_API_EXPORT iree_status_t IREE_API_CALL @@ -2171,7 +2170,7 @@ IREE_TRACE_SCOPE0("iree_hal_semaphore_wait_with_timeout"); auto* handle = reinterpret_cast<Semaphore*>(semaphore); if (!handle) return IREE_STATUS_INVALID_ARGUMENT; - return ToApiStatus(handle->Wait(value, ToAbslDuration(timeout_ns))); + return ToApiStatus(handle->Wait(value, Duration(timeout_ns))); } } // namespace hal
diff --git a/iree/hal/command_queue.h b/iree/hal/command_queue.h index 8e82e8f..7d068d7 100644 --- a/iree/hal/command_queue.h +++ b/iree/hal/command_queue.h
@@ -18,8 +18,6 @@ #include <cstdint> #include <string> -#include "absl/time/clock.h" -#include "absl/time/time.h" #include "absl/types/span.h" #include "iree/base/bitfield.h" #include "iree/base/status.h" @@ -91,11 +89,11 @@ // // If the command queue has encountered an error during submission at any // point it will be returned here (repeatedly). - virtual Status WaitIdle(absl::Time deadline) = 0; - inline Status WaitIdle(absl::Duration timeout) { - return WaitIdle(RelativeTimeoutToDeadline(timeout)); + virtual Status WaitIdle(Time deadline_ns) = 0; + inline Status WaitIdle(Duration timeout_ns) { + return WaitIdle(RelativeTimeoutToDeadlineNanos(timeout_ns)); } - inline Status WaitIdle() { return WaitIdle(absl::InfiniteFuture()); } + inline Status WaitIdle() { return WaitIdle(InfiniteFuture()); } protected: CommandQueue(std::string name, CommandCategoryBitfield supported_categories)
diff --git a/iree/hal/cts/command_queue_test.cc b/iree/hal/cts/command_queue_test.cc index 94ce5d3..adae1df 100644 --- a/iree/hal/cts/command_queue_test.cc +++ b/iree/hal/cts/command_queue_test.cc
@@ -71,7 +71,7 @@ ASSERT_OK(command_queue->Submit( {{}, {command_buffer.get()}, {{semaphore.get(), 1ull}}})); - ASSERT_OK(semaphore->Wait(1ull, absl::InfiniteFuture())); + ASSERT_OK(semaphore->Wait(1ull, InfiniteFuture())); } // Tests waiting while work is pending/in-flight. @@ -91,12 +91,11 @@ // Work shouldn't start until the wait semaphore reaches its payload value. EXPECT_THAT(signal_semaphore->Query(), IsOkAndHolds(Eq(0ull))); - EXPECT_TRUE( - IsDeadlineExceeded(command_queue->WaitIdle(absl::Milliseconds(100)))); + EXPECT_TRUE(IsDeadlineExceeded(command_queue->WaitIdle(Milliseconds(100)))); // Signal the wait semaphore, work should begin and complete. ASSERT_OK(wait_semaphore->Signal(1ull)); - ASSERT_OK(signal_semaphore->Wait(1ull, absl::InfiniteFuture())); + ASSERT_OK(signal_semaphore->Wait(1ull, InfiniteFuture())); } // Tests using multiple wait and signal semaphores. @@ -120,8 +119,7 @@ EXPECT_THAT(signal_semaphore_1->Query(), IsOkAndHolds(Eq(0ull))); EXPECT_THAT(signal_semaphore_2->Query(), IsOkAndHolds(Eq(0ull))); // Note: This fails with Vulkan timeline semaphore emulation (returns OK) - EXPECT_TRUE( - IsDeadlineExceeded(command_queue->WaitIdle(absl::Milliseconds(100)))); + EXPECT_TRUE(IsDeadlineExceeded(command_queue->WaitIdle(Milliseconds(100)))); // Signal the wait semaphores, work should only begin after each is set. ASSERT_OK(wait_semaphore_1->Signal(1ull));
diff --git a/iree/hal/cts/semaphore_test.cc b/iree/hal/cts/semaphore_test.cc index 873837c..5b212aa 100644 --- a/iree/hal/cts/semaphore_test.cc +++ b/iree/hal/cts/semaphore_test.cc
@@ -12,6 +12,8 @@ // See the License for the specific language governing permissions and // limitations under the License. +#include <thread> + #include "iree/hal/cts/cts_test_base.h" #include "iree/hal/driver_registry.h" #include "iree/testing/gtest.h" @@ -60,17 +62,17 @@ // Tests waiting on no semaphores. TEST_P(SemaphoreTest, EmptyWait) { - EXPECT_OK(device_->WaitAllSemaphores({}, absl::InfiniteFuture())); + EXPECT_OK(device_->WaitAllSemaphores({}, InfiniteFuture())); } // Tests waiting on a semaphore that has already been signaled. TEST_P(SemaphoreTest, WaitAlreadySignaled) { ASSERT_OK_AND_ASSIGN(auto semaphore, device_->CreateSemaphore(2u)); // Test both previous and current values. - EXPECT_OK(device_->WaitAllSemaphores({{semaphore.get(), 1u}}, - absl::InfiniteFuture())); - EXPECT_OK(device_->WaitAllSemaphores({{semaphore.get(), 2u}}, - absl::InfiniteFuture())); + EXPECT_OK( + device_->WaitAllSemaphores({{semaphore.get(), 1u}}, InfiniteFuture())); + EXPECT_OK( + device_->WaitAllSemaphores({{semaphore.get(), 2u}}, InfiniteFuture())); } // Tests waiting on a semaphore that has not been signaled. @@ -79,7 +81,7 @@ // NOTE: we don't actually block here because otherwise we'd lock up. // Result status is undefined - some backends may return DeadlineExceededError // while others may return success. - device_->WaitAllSemaphores({{semaphore.get(), 3u}}, absl::InfinitePast()) + device_->WaitAllSemaphores({{semaphore.get(), 3u}}, InfinitePast()) .IgnoreError(); } @@ -93,15 +95,12 @@ ASSERT_OK_AND_ASSIGN(auto b2a, device_->CreateSemaphore(0u)); std::thread thread([&]() { // Should advance right past this because the value is already set. - ASSERT_OK( - device_->WaitAllSemaphores({{a2b.get(), 0u}}, absl::InfiniteFuture())); + ASSERT_OK(device_->WaitAllSemaphores({{a2b.get(), 0u}}, InfiniteFuture())); ASSERT_OK(b2a->Signal(1u)); // Jump ahead. - ASSERT_OK( - device_->WaitAllSemaphores({{a2b.get(), 4u}}, absl::InfiniteFuture())); + ASSERT_OK(device_->WaitAllSemaphores({{a2b.get(), 4u}}, InfiniteFuture())); }); - ASSERT_OK( - device_->WaitAllSemaphores({{b2a.get(), 1u}}, absl::InfiniteFuture())); + ASSERT_OK(device_->WaitAllSemaphores({{b2a.get(), 1u}}, InfiniteFuture())); ASSERT_OK(a2b->Signal(4u)); thread.join(); }
diff --git a/iree/hal/device.h b/iree/hal/device.h index 9bcf1db..4cd3696 100644 --- a/iree/hal/device.h +++ b/iree/hal/device.h
@@ -17,8 +17,6 @@ #include <memory> -#include "absl/time/clock.h" -#include "absl/time/time.h" #include "iree/base/ref_ptr.h" #include "iree/base/status.h" #include "iree/base/target_platform.h" @@ -138,10 +136,11 @@ // having been signaled. Note that a subset of the |semaphores| may have been // signaled and each can be queried to see which ones. virtual Status WaitAllSemaphores(absl::Span<const SemaphoreValue> semaphores, - absl::Time deadline) = 0; + Time deadline_ns) = 0; inline Status WaitAllSemaphores(absl::Span<const SemaphoreValue> semaphores, - absl::Duration timeout) { - return WaitAllSemaphores(semaphores, RelativeTimeoutToDeadline(timeout)); + Duration timeout_ns) { + return WaitAllSemaphores(semaphores, + RelativeTimeoutToDeadlineNanos(timeout_ns)); } // Blocks the caller until at least one of the |semaphores| reaches or exceeds @@ -156,20 +155,21 @@ // Returns DEADLINE_EXCEEDED if the |deadline| elapses without any semaphores // having been signaled. virtual StatusOr<int> WaitAnySemaphore( - absl::Span<const SemaphoreValue> semaphores, absl::Time deadline) = 0; + absl::Span<const SemaphoreValue> semaphores, Time deadline_ns) = 0; inline StatusOr<int> WaitAnySemaphore( - absl::Span<const SemaphoreValue> semaphores, absl::Duration timeout) { - return WaitAnySemaphore(semaphores, RelativeTimeoutToDeadline(timeout)); + absl::Span<const SemaphoreValue> semaphores, Duration timeout_ns) { + return WaitAnySemaphore(semaphores, + RelativeTimeoutToDeadlineNanos(timeout_ns)); } // Blocks until all outstanding requests on all queues have been // completed. This is equivalent to having waited on all outstanding // semaphores. - virtual Status WaitIdle(absl::Time deadline) = 0; - inline Status WaitIdle(absl::Duration timeout) { - return WaitIdle(RelativeTimeoutToDeadline(timeout)); + virtual Status WaitIdle(Time deadline_ns) = 0; + inline Status WaitIdle(Duration timeout_ns) { + return WaitIdle(RelativeTimeoutToDeadlineNanos(timeout_ns)); } - inline Status WaitIdle() { return WaitIdle(absl::InfiniteFuture()); } + inline Status WaitIdle() { return WaitIdle(InfiniteFuture()); } protected: explicit Device(DeviceInfo device_info)
diff --git a/iree/hal/device_manager.cc b/iree/hal/device_manager.cc index 2646204..30f3c80 100644 --- a/iree/hal/device_manager.cc +++ b/iree/hal/device_manager.cc
@@ -178,7 +178,7 @@ Status DeviceManager::Submit(Device* device, CommandQueue* command_queue, absl::Span<const SubmissionBatch> batches, - absl::Time deadline) { + Time deadline_ns) { IREE_TRACE_SCOPE0("DeviceManager::Submit"); return command_queue->Submit(batches); } @@ -188,11 +188,11 @@ return OkStatus(); } -Status DeviceManager::WaitIdle(absl::Time deadline) { +Status DeviceManager::WaitIdle(Time deadline_ns) { IREE_TRACE_SCOPE0("DeviceManager::WaitIdle"); absl::MutexLock lock(&device_mutex_); for (const auto& device : devices_) { - RETURN_IF_ERROR(device->WaitIdle(deadline)); + RETURN_IF_ERROR(device->WaitIdle(deadline_ns)); } return OkStatus(); }
diff --git a/iree/hal/device_manager.h b/iree/hal/device_manager.h index faaddd9..492b824 100644 --- a/iree/hal/device_manager.h +++ b/iree/hal/device_manager.h
@@ -18,8 +18,6 @@ #include <vector> #include "absl/synchronization/mutex.h" -#include "absl/time/clock.h" -#include "absl/time/time.h" #include "absl/types/span.h" #include "iree/base/status.h" #include "iree/base/time.h" @@ -164,16 +162,16 @@ // Submissions may be made from any thread. Behavior is undefined // if a thread is performing a WaitIdle while another thread submits work. Status Submit(Device* device, CommandQueue* command_queue, - absl::Span<const SubmissionBatch> batches, absl::Time deadline); + absl::Span<const SubmissionBatch> batches, Time deadline_ns); Status Submit(Device* device, CommandQueue* command_queue, absl::Span<const SubmissionBatch> batches, - absl::Duration timeout) { + Duration timeout_ns) { return Submit(device, command_queue, batches, - RelativeTimeoutToDeadline(timeout)); + RelativeTimeoutToDeadlineNanos(timeout_ns)); } Status Submit(Device* device, CommandQueue* command_queue, absl::Span<const SubmissionBatch> batches) { - return Submit(device, command_queue, batches, absl::InfinitePast()); + return Submit(device, command_queue, batches, InfinitePast()); } // Flushes any requests that are pending in the scheduler and ensures they @@ -192,11 +190,11 @@ // // If any used device has encountered an error during submission at any // point it will be returned here (repeatedly). - Status WaitIdle(absl::Time deadline); - inline Status WaitIdle(absl::Duration timeout) { - return WaitIdle(RelativeTimeoutToDeadline(timeout)); + Status WaitIdle(Time deadline_ns); + inline Status WaitIdle(Duration timeout_ns) { + return WaitIdle(RelativeTimeoutToDeadlineNanos(timeout_ns)); } - inline Status WaitIdle() { return WaitIdle(absl::InfiniteFuture()); } + inline Status WaitIdle() { return WaitIdle(InfiniteFuture()); } private: mutable absl::Mutex device_mutex_;
diff --git a/iree/hal/dylib/BUILD b/iree/hal/dylib/BUILD index 25c08ea..fc3ccb9 100644 --- a/iree/hal/dylib/BUILD +++ b/iree/hal/dylib/BUILD
@@ -60,7 +60,6 @@ srcs = ["dylib_executable.cc"], hdrs = ["dylib_executable.h"], deps = [ - ":memref_runtime", "//iree/base:dynamic_library", "//iree/base:file_io", "//iree/base:status", @@ -89,10 +88,3 @@ "//iree/hal:executable_format", ], ) - -cc_library( - name = "memref_runtime", - hdrs = [ - "memref_runtime.h", - ], -)
diff --git a/iree/hal/dylib/CMakeLists.txt b/iree/hal/dylib/CMakeLists.txt index 7644d92..d720435 100644 --- a/iree/hal/dylib/CMakeLists.txt +++ b/iree/hal/dylib/CMakeLists.txt
@@ -65,7 +65,6 @@ SRCS "dylib_executable.cc" DEPS - ::memref_runtime absl::inlined_vector absl::span flatbuffers @@ -97,11 +96,3 @@ iree::hal::executable_format PUBLIC ) - -iree_cc_library( - NAME - memref_runtime - HDRS - "memref_runtime.h" - PUBLIC -)
diff --git a/iree/hal/dylib/dylib_executable.cc b/iree/hal/dylib/dylib_executable.cc index e06bb19..e58a003 100644 --- a/iree/hal/dylib/dylib_executable.cc +++ b/iree/hal/dylib/dylib_executable.cc
@@ -17,7 +17,6 @@ #include "flatbuffers/flatbuffers.h" #include "iree/base/file_io.h" #include "iree/base/tracing.h" -#include "iree/hal/dylib/memref_runtime.h" #include "iree/schemas/dylib_executable_def_generated.h" namespace iree { @@ -96,15 +95,9 @@ struct DyLibDispatchState : public HostExecutable::DispatchState { DyLibDispatchState() = default; - ~DyLibDispatchState() override { - for (int i = 0; i < descriptors.size(); ++i) { - freeUnrankedDescriptor(descriptors[i]); - } - } - void* entry_function = nullptr; - absl::InlinedVector<UnrankedMemRefType<uint32_t>*, 4> descriptors; absl::InlinedVector<void*, 4> args; + absl::InlinedVector<int64_t, 4> push_constant; }; StatusOr<ref_ptr<HostExecutable::DispatchState>> @@ -127,17 +120,14 @@ MemoryAccessBitfield::kWrite, io_binding.offset, io_binding.length)); auto data = memory.mutable_data(); - auto descriptor = allocUnrankedDescriptor<uint32_t>(data); - dispatch_state->descriptors.push_back(descriptor); - dispatch_state->args.push_back(&descriptor->descriptor); + + dispatch_state->args.push_back(data); } } - - auto push_constants_descriptor = allocUnrankedDescriptor<uint32_t>( - const_cast<uint32_t*>(params.push_constants->values.data()), - {static_cast<int64_t>(params.push_constants->values.size())}); - dispatch_state->descriptors.push_back(push_constants_descriptor); - dispatch_state->args.push_back(&push_constants_descriptor->descriptor); + // TODO(ataei): Consider moving this casting to codegen side ?! + for (int i = 0; i < params.push_constants->values.size(); ++i) { + dispatch_state->push_constant.push_back(params.push_constants->values[i]); + } return std::move(dispatch_state); } @@ -147,8 +137,10 @@ IREE_TRACE_SCOPE0("DyLibExecutable::DispatchTile"); auto* dispatch_state = static_cast<DyLibDispatchState*>(state); - auto entry_function = (void (*)(void**))dispatch_state->entry_function; - entry_function(dispatch_state->args.data()); + auto entry_function = + (void (*)(void**, int64_t*))dispatch_state->entry_function; + entry_function(dispatch_state->args.data(), + dispatch_state->push_constant.data()); return OkStatus(); }
diff --git a/iree/hal/dylib/memref_runtime.h b/iree/hal/dylib/memref_runtime.h deleted file mode 100644 index 50d3987..0000000 --- a/iree/hal/dylib/memref_runtime.h +++ /dev/null
@@ -1,177 +0,0 @@ -// Copyright 2020 Google LLC -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// https://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// - -#ifndef IREE_HAL_DYLIB_MEMREF_RUNTIME_H_ -#define IREE_HAL_DYLIB_MEMREF_RUNTIME_H_ - -#include <assert.h> - -#include <cstdint> -#include <vector> - -namespace iree { -namespace hal { -namespace dylib { - -template <int N> -void dropFront(int64_t arr[N], int64_t *res) { - for (unsigned i = 1; i < N; ++i) *(res + i - 1) = arr[i]; -} - -/// StridedMemRef descriptor type with static rank. -template <typename T, int N> -struct StridedMemRefType { - T *basePtr; - T *data; - int64_t offset; - int64_t sizes[N]; - int64_t strides[N]; - // This operator[] is extremely slow and only for sugaring purposes. - StridedMemRefType<T, N - 1> operator[](int64_t idx) { - StridedMemRefType<T, N - 1> res; - res.basePtr = basePtr; - res.data = data; - res.offset = offset + idx * strides[0]; - dropFront<N>(sizes, res.sizes); - dropFront<N>(strides, res.strides); - return res; - } -}; - -/// StridedMemRef descriptor type specialized for rank 1. -template <typename T> -struct StridedMemRefType<T, 1> { - T *basePtr; - T *data; - int64_t offset; - int64_t sizes[1]; - int64_t strides[1]; - T &operator[](int64_t idx) { return *(data + offset + idx * strides[0]); } -}; - -/// StridedMemRef descriptor type specialized for rank 0. -template <typename T> -struct StridedMemRefType<T, 0> { - T *basePtr; - T *data; - int64_t offset; -}; - -// Unranked MemRef -template <typename T> -struct UnrankedMemRefType { - int64_t rank; - void *descriptor; -}; - -// Given a shape with sizes greater than 0 along all dimensions, -// returns the distance, in number of elements, between a slice in a dimension -// and the next slice in the same dimension. -// e.g. shape[3, 4, 5] -> strides[20, 5, 1] -inline std::vector<int64_t> makeStrides(const std::vector<int64_t> &shape) { - std::vector<int64_t> tmp; - if (shape.empty()) return tmp; - tmp.reserve(shape.size()); - int64_t running = 1; - for (auto rit = shape.rbegin(), reit = shape.rend(); rit != reit; ++rit) { - assert(*rit > 0 && - "size must be greater than 0 along all dimensions of shape"); - tmp.push_back(running); - running *= *rit; - } - return std::vector<int64_t>(tmp.rbegin(), tmp.rend()); -} - -// Mallocs a StridedMemRefDescriptor<T, N>* that matches the MLIR ABI. -// This is an implementation detail that is kept in sync with MLIR codegen -// conventions. -template <typename T, int N> -StridedMemRefType<T, N> *makeStridedMemRefDescriptor( - void *ptr, const std::vector<int64_t> &shape) { - StridedMemRefType<T, N> *descriptor = static_cast<StridedMemRefType<T, N> *>( - malloc(sizeof(StridedMemRefType<T, N>))); - descriptor->basePtr = static_cast<T *>(ptr); - descriptor->data = static_cast<T *>(ptr); - descriptor->offset = 0; - std::copy(shape.begin(), shape.end(), descriptor->sizes); - auto strides = makeStrides(shape); - std::copy(strides.begin(), strides.end(), descriptor->strides); - return descriptor; -} - -// Mallocs a StridedMemRefDescriptor<T, 0>* (i.e. a pointer to scalar) that -// matches the MLIR ABI. This is an implementation detail that is kept in sync -// with MLIR codegen conventions. -template <typename T> -StridedMemRefType<T, 0> *makeStridedMemRefDescriptor( - void *ptr, const std::vector<int64_t> &shape) { - StridedMemRefType<T, 0> *descriptor = static_cast<StridedMemRefType<T, 0> *>( - malloc(sizeof(StridedMemRefType<T, 0>))); - descriptor->basePtr = static_cast<T *>(ptr); - descriptor->data = static_cast<T *>(ptr); - descriptor->offset = 0; - return descriptor; -} - -// Mallocs an UnrankedMemRefType<T>* that contains a ranked -// StridedMemRefDescriptor<T, Rank>* and matches the MLIR ABI. This is an -// implementation detail that is kept in sync with MLIR codegen conventions. -template <typename T> -UnrankedMemRefType<T> *allocUnrankedDescriptor( - void *data, const std::vector<int64_t> &shape) { - UnrankedMemRefType<T> *res = static_cast<UnrankedMemRefType<T> *>( - malloc(sizeof(UnrankedMemRefType<T>))); - res->rank = shape.size(); - if (res->rank == 0) - res->descriptor = makeStridedMemRefDescriptor<T>(data, shape); - else if (res->rank == 1) - res->descriptor = makeStridedMemRefDescriptor<T, 1>(data, shape); - else if (res->rank == 2) - res->descriptor = makeStridedMemRefDescriptor<T, 2>(data, shape); - else if (res->rank == 3) - res->descriptor = makeStridedMemRefDescriptor<T, 3>(data, shape); - else if (res->rank == 4) - res->descriptor = makeStridedMemRefDescriptor<T, 4>(data, shape); - else if (res->rank == 5) - res->descriptor = makeStridedMemRefDescriptor<T, 5>(data, shape); - else if (res->rank == 6) - res->descriptor = makeStridedMemRefDescriptor<T, 6>(data, shape); - else - assert(false && "Unsupported 6+D memref descriptor"); - return res; -} - -// Shape and strides aren't used in the generated code (yet). -// TODO(ataei): Delete this version once we can pass shapes. -template <typename T> -UnrankedMemRefType<T> *allocUnrankedDescriptor(void *data) { - UnrankedMemRefType<T> *res = static_cast<UnrankedMemRefType<T> *>( - malloc(sizeof(UnrankedMemRefType<T>))); - res->descriptor = makeStridedMemRefDescriptor<T>(data, {}); - return res; -} - -// Frees an UnrankedMemRefType<T>* -template <typename T> -void freeUnrankedDescriptor(UnrankedMemRefType<T> *desc) { - free(desc->descriptor); - free(desc); -} - -} // namespace dylib -} // namespace hal -} // namespace iree - -#endif // IREE_HAL_DYLIB_MEMREF_RUNTIME_H_
diff --git a/iree/hal/host/BUILD b/iree/hal/host/BUILD index 3480cc4..9e25ffa 100644 --- a/iree/hal/host/BUILD +++ b/iree/hal/host/BUILD
@@ -41,10 +41,10 @@ srcs = ["condvar_semaphore_test.cc"], deps = [ ":condvar_semaphore", + "//iree/base:api", "//iree/base:status", "//iree/base:status_matchers", "//iree/testing:gtest_main", - "@com_google_absl//absl/time", ], )
diff --git a/iree/hal/host/CMakeLists.txt b/iree/hal/host/CMakeLists.txt index 06383f3..8b01d00 100644 --- a/iree/hal/host/CMakeLists.txt +++ b/iree/hal/host/CMakeLists.txt
@@ -39,7 +39,7 @@ "condvar_semaphore_test.cc" DEPS ::condvar_semaphore - absl::time + iree::base::api iree::base::status iree::base::status_matchers iree::testing::gtest_main
diff --git a/iree/hal/host/condvar_semaphore.cc b/iree/hal/host/condvar_semaphore.cc index 2a3fb0c..1f79a25 100644 --- a/iree/hal/host/condvar_semaphore.cc +++ b/iree/hal/host/condvar_semaphore.cc
@@ -60,7 +60,7 @@ // static Status CondVarSemaphore::WaitForSemaphores( absl::Span<const SemaphoreValue> semaphores, bool wait_all, - absl::Time deadline) { + Time deadline_ns) { IREE_TRACE_SCOPE0("CondVarSemaphore::WaitForSemaphores"); // Some of the semaphores may already be signaled; we only need to wait for @@ -94,7 +94,7 @@ semaphore_value->second; }, &semaphore_value), - deadline)) { + absl::FromUnixNanos(static_cast<int64_t>(deadline_ns)))) { return DeadlineExceededErrorBuilder(IREE_LOC) << "Deadline exceeded waiting for semaphores"; } @@ -106,8 +106,8 @@ return OkStatus(); } -Status CondVarSemaphore::Wait(uint64_t value, absl::Time deadline) { - return WaitForSemaphores({{this, value}}, /*wait_all=*/true, deadline); +Status CondVarSemaphore::Wait(uint64_t value, Time deadline_ns) { + return WaitForSemaphores({{this, value}}, /*wait_all=*/true, deadline_ns); } } // namespace host
diff --git a/iree/hal/host/condvar_semaphore.h b/iree/hal/host/condvar_semaphore.h index daf3825..44dd541 100644 --- a/iree/hal/host/condvar_semaphore.h +++ b/iree/hal/host/condvar_semaphore.h
@@ -37,7 +37,7 @@ // Waits for one or more (or all) semaphores to reach or exceed the given // values. static Status WaitForSemaphores(absl::Span<const SemaphoreValue> semaphores, - bool wait_all, absl::Time deadline); + bool wait_all, Time deadline_ns); explicit CondVarSemaphore(uint64_t initial_value); ~CondVarSemaphore() override; @@ -46,7 +46,7 @@ Status Signal(uint64_t value) override; void Fail(Status status) override; - Status Wait(uint64_t value, absl::Time deadline) override; + Status Wait(uint64_t value, Time deadline_ns) override; private: // The mutex is not required to query the value; this lets us quickly check if
diff --git a/iree/hal/host/condvar_semaphore_test.cc b/iree/hal/host/condvar_semaphore_test.cc index c3614f5..9846f29 100644 --- a/iree/hal/host/condvar_semaphore_test.cc +++ b/iree/hal/host/condvar_semaphore_test.cc
@@ -17,7 +17,6 @@ #include <cstdint> #include <thread> // NOLINT -#include "absl/time/time.h" #include "iree/base/status.h" #include "iree/base/status_matchers.h" #include "iree/testing/gtest.h" @@ -73,7 +72,7 @@ // Tests waiting on no semaphores. TEST(CondVarSemaphoreTest, EmptyWait) { EXPECT_OK(CondVarSemaphore::WaitForSemaphores({}, /*wait_all=*/true, - absl::InfiniteFuture())); + InfiniteFuture())); } // Tests waiting on a semaphore that has already been signaled. @@ -81,9 +80,9 @@ CondVarSemaphore semaphore(2u); // Test both previous and current values. EXPECT_OK(CondVarSemaphore::WaitForSemaphores( - {{&semaphore, 1u}}, /*wait_all=*/true, absl::InfiniteFuture())); + {{&semaphore, 1u}}, /*wait_all=*/true, InfiniteFuture())); EXPECT_OK(CondVarSemaphore::WaitForSemaphores( - {{&semaphore, 2u}}, /*wait_all=*/true, absl::InfiniteFuture())); + {{&semaphore, 2u}}, /*wait_all=*/true, InfiniteFuture())); } // Tests waiting on a semaphore that has not been signaled. @@ -91,7 +90,7 @@ CondVarSemaphore semaphore(2u); // NOTE: we don't actually block here because otherwise we'd lock up. EXPECT_TRUE(IsDeadlineExceeded(CondVarSemaphore::WaitForSemaphores( - {{&semaphore, 3u}}, /*wait_all=*/true, absl::InfinitePast()))); + {{&semaphore, 3u}}, /*wait_all=*/true, InfinitePast()))); } // Tests waiting on a failed semaphore (it should return the error on the @@ -100,7 +99,7 @@ CondVarSemaphore semaphore(2u); semaphore.Fail(UnknownErrorBuilder(IREE_LOC)); EXPECT_TRUE(IsUnknown(CondVarSemaphore::WaitForSemaphores( - {{&semaphore, 2u}}, /*wait_all=*/true, absl::InfinitePast()))); + {{&semaphore, 2u}}, /*wait_all=*/true, InfinitePast()))); } // Tests threading behavior by ping-ponging between the test main thread and @@ -111,14 +110,14 @@ std::thread thread([&]() { // Should advance right past this because the value is already set. ASSERT_OK(CondVarSemaphore::WaitForSemaphores( - {{&a2b, 0u}}, /*wait_all=*/true, absl::InfiniteFuture())); + {{&a2b, 0u}}, /*wait_all=*/true, InfiniteFuture())); ASSERT_OK(b2a.Signal(1u)); // Jump ahead. ASSERT_OK(CondVarSemaphore::WaitForSemaphores( - {{&a2b, 4u}}, /*wait_all=*/true, absl::InfiniteFuture())); + {{&a2b, 4u}}, /*wait_all=*/true, InfiniteFuture())); }); ASSERT_OK(CondVarSemaphore::WaitForSemaphores({{&b2a, 1u}}, /*wait_all=*/true, - absl::InfiniteFuture())); + InfiniteFuture())); ASSERT_OK(a2b.Signal(4u)); thread.join(); } @@ -131,10 +130,10 @@ std::thread thread([&]() { ASSERT_OK(b2a.Signal(1u)); got_failure = IsUnknown(CondVarSemaphore::WaitForSemaphores( - {{&a2b, 1u}}, /*wait_all=*/true, absl::InfiniteFuture())); + {{&a2b, 1u}}, /*wait_all=*/true, InfiniteFuture())); }); ASSERT_OK(CondVarSemaphore::WaitForSemaphores({{&b2a, 1u}}, /*wait_all=*/true, - absl::InfiniteFuture())); + InfiniteFuture())); a2b.Fail(UnknownErrorBuilder(IREE_LOC)); thread.join(); ASSERT_TRUE(got_failure);
diff --git a/iree/hal/host/host_local_device.cc b/iree/hal/host/host_local_device.cc index 5638b9a..1ca3c54 100644 --- a/iree/hal/host/host_local_device.cc +++ b/iree/hal/host/host_local_device.cc
@@ -77,20 +77,20 @@ } Status HostLocalDevice::WaitAllSemaphores( - absl::Span<const SemaphoreValue> semaphores, absl::Time deadline) { + absl::Span<const SemaphoreValue> semaphores, Time deadline_ns) { IREE_TRACE_SCOPE0("HostLocalDevice::WaitAllSemaphores"); - return scheduling_model_->WaitAllSemaphores(semaphores, deadline); + return scheduling_model_->WaitAllSemaphores(semaphores, deadline_ns); } StatusOr<int> HostLocalDevice::WaitAnySemaphore( - absl::Span<const SemaphoreValue> semaphores, absl::Time deadline) { + absl::Span<const SemaphoreValue> semaphores, Time deadline_ns) { IREE_TRACE_SCOPE0("HostLocalDevice::WaitAnySemaphore"); - return scheduling_model_->WaitAnySemaphore(semaphores, deadline); + return scheduling_model_->WaitAnySemaphore(semaphores, deadline_ns); } -Status HostLocalDevice::WaitIdle(absl::Time deadline) { +Status HostLocalDevice::WaitIdle(Time deadline_ns) { IREE_TRACE_SCOPE0("HostLocalDevice::WaitIdle"); - return scheduling_model_->WaitIdle(deadline); + return scheduling_model_->WaitIdle(deadline_ns); } } // namespace host
diff --git a/iree/hal/host/host_local_device.h b/iree/hal/host/host_local_device.h index de00d5f..8fc3b31 100644 --- a/iree/hal/host/host_local_device.h +++ b/iree/hal/host/host_local_device.h
@@ -63,11 +63,11 @@ StatusOr<ref_ptr<Semaphore>> CreateSemaphore(uint64_t initial_value) override; Status WaitAllSemaphores(absl::Span<const SemaphoreValue> semaphores, - absl::Time deadline) override; + Time deadline_ns) override; StatusOr<int> WaitAnySemaphore(absl::Span<const SemaphoreValue> semaphores, - absl::Time deadline) override; + Time deadline_ns) override; - Status WaitIdle(absl::Time deadline) override; + Status WaitIdle(Time deadline_ns) override; protected: explicit HostLocalDevice(DeviceInfo device_info,
diff --git a/iree/hal/host/scheduling_model.h b/iree/hal/host/scheduling_model.h index df56b51..38771ec 100644 --- a/iree/hal/host/scheduling_model.h +++ b/iree/hal/host/scheduling_model.h
@@ -71,7 +71,7 @@ // having been signaled. Note that a subset of the |semaphores| may have been // signaled and each can be queried to see which ones. virtual Status WaitAllSemaphores(absl::Span<const SemaphoreValue> semaphores, - absl::Time deadline) = 0; + Time deadline_ns) = 0; // Blocks the caller until at least one of the |semaphores| reaches or exceeds // the specified payload value or the |deadline| elapses. All |semaphores| @@ -85,12 +85,12 @@ // Returns DEADLINE_EXCEEDED if the |deadline| elapses without any semaphores // having been signaled. virtual StatusOr<int> WaitAnySemaphore( - absl::Span<const SemaphoreValue> semaphores, absl::Time deadline) = 0; + absl::Span<const SemaphoreValue> semaphores, Time deadline_ns) = 0; // Blocks until all outstanding requests on all queues have been // completed. This is equivalent to having waited on all outstanding // semaphores. - virtual Status WaitIdle(absl::Time deadline) = 0; + virtual Status WaitIdle(Time deadline_ns) = 0; }; } // namespace host
diff --git a/iree/hal/host/serial/BUILD b/iree/hal/host/serial/BUILD index 04281e4..b99ae45 100644 --- a/iree/hal/host/serial/BUILD +++ b/iree/hal/host/serial/BUILD
@@ -50,7 +50,6 @@ "//iree/hal/testing:mock_command_queue", "//iree/testing:gtest_main", "@com_google_absl//absl/memory", - "@com_google_absl//absl/time", ], )
diff --git a/iree/hal/host/serial/CMakeLists.txt b/iree/hal/host/serial/CMakeLists.txt index 93b5c73..954f988 100644 --- a/iree/hal/host/serial/CMakeLists.txt +++ b/iree/hal/host/serial/CMakeLists.txt
@@ -40,7 +40,6 @@ DEPS ::async_command_queue absl::memory - absl::time iree::base::status iree::base::status_matchers iree::base::time
diff --git a/iree/hal/host/serial/async_command_queue.cc b/iree/hal/host/serial/async_command_queue.cc index 8b789ca..e1b8fac 100644 --- a/iree/hal/host/serial/async_command_queue.cc +++ b/iree/hal/host/serial/async_command_queue.cc
@@ -103,7 +103,7 @@ return submission_queue_.Enqueue(batches); } -Status AsyncCommandQueue::WaitIdle(absl::Time deadline) { +Status AsyncCommandQueue::WaitIdle(Time deadline_ns) { IREE_TRACE_SCOPE0("AsyncCommandQueue::WaitIdle"); // Wait until the deadline, the thread exits, or there are no more pending @@ -115,7 +115,7 @@ return queue->empty() || !queue->permanent_error().ok(); }, &submission_queue_), - deadline)) { + absl::FromUnixNanos(static_cast<int64_t>(deadline_ns)))) { return DeadlineExceededErrorBuilder(IREE_LOC) << "Deadline exceeded waiting for submission thread to go idle"; }
diff --git a/iree/hal/host/serial/async_command_queue.h b/iree/hal/host/serial/async_command_queue.h index 76d4797..e474c23 100644 --- a/iree/hal/host/serial/async_command_queue.h +++ b/iree/hal/host/serial/async_command_queue.h
@@ -47,7 +47,7 @@ Status Submit(absl::Span<const SubmissionBatch> batches) override; - Status WaitIdle(absl::Time deadline) override; + Status WaitIdle(Time deadline_ns) override; private: // Thread entry point for the async worker thread.
diff --git a/iree/hal/host/serial/async_command_queue_test.cc b/iree/hal/host/serial/async_command_queue_test.cc index 06ca4ae..5aca399 100644 --- a/iree/hal/host/serial/async_command_queue_test.cc +++ b/iree/hal/host/serial/async_command_queue_test.cc
@@ -14,13 +14,13 @@ #include "iree/hal/host/serial/async_command_queue.h" +#include <chrono> #include <cstdint> #include <memory> +#include <thread> #include <utility> #include "absl/memory/memory.h" -#include "absl/time/clock.h" -#include "absl/time/time.h" #include "iree/base/status.h" #include "iree/base/status_matchers.h" #include "iree/base/time.h" @@ -40,6 +40,13 @@ using testing::MockCommandBuffer; using testing::MockCommandQueue; +// Suspends execution of the calling thread for the given |duration_ms|. +// Depending on platform this may have an extremely coarse resolution (upwards +// of several to dozens of milliseconds). +inline void Sleep(std::chrono::milliseconds duration_ms) { + std::this_thread::sleep_for(duration_ms); +} + struct AsyncCommandQueueTest : public ::testing::Test { MockCommandQueue* mock_target_queue; std::unique_ptr<CommandQueue> command_queue; @@ -75,7 +82,7 @@ CondVarSemaphore semaphore(0ull); ASSERT_OK( command_queue->Submit({{}, {cmd_buffer.get()}, {{&semaphore, 1ull}}})); - ASSERT_OK(semaphore.Wait(1ull, absl::InfiniteFuture())); + ASSERT_OK(semaphore.Wait(1ull, InfiniteFuture())); } // Tests that failure is propagated along the fence from the target queue. @@ -92,7 +99,7 @@ CondVarSemaphore semaphore(0ull); ASSERT_OK( command_queue->Submit({{}, {cmd_buffer.get()}, {{&semaphore, 1ull}}})); - EXPECT_TRUE(IsDataLoss(semaphore.Wait(1ull, absl::InfiniteFuture()))); + EXPECT_TRUE(IsDataLoss(semaphore.Wait(1ull, InfiniteFuture()))); } // Tests that waiting for idle is a no-op when nothing is queued. @@ -109,7 +116,7 @@ EXPECT_CALL(*mock_target_queue, Submit(_)) .WillOnce([](absl::Span<const SubmissionBatch> batches) { - Sleep(absl::Milliseconds(100)); + Sleep(std::chrono::milliseconds(100)); return OkStatus(); }); CondVarSemaphore semaphore(0ull); @@ -131,7 +138,7 @@ EXPECT_CALL(*mock_target_queue, Submit(_)) .WillRepeatedly([](absl::Span<const SubmissionBatch> batches) { - Sleep(absl::Milliseconds(100)); + Sleep(std::chrono::milliseconds(100)); return OkStatus(); }); @@ -164,7 +171,7 @@ // Fail. EXPECT_CALL(*mock_target_queue, Submit(_)) .WillOnce([](absl::Span<const SubmissionBatch> batches) { - Sleep(absl::Milliseconds(100)); + Sleep(std::chrono::milliseconds(100)); return DataLossErrorBuilder(IREE_LOC); }); auto cmd_buffer_0 = make_ref<MockCommandBuffer>(CommandBufferMode::kOneShot, @@ -172,7 +179,7 @@ CondVarSemaphore semaphore_0(0ull); ASSERT_OK( command_queue->Submit({{}, {cmd_buffer_0.get()}, {{&semaphore_0, 1u}}})); - EXPECT_TRUE(IsDataLoss(semaphore_0.Wait(1ull, absl::InfiniteFuture()))); + EXPECT_TRUE(IsDataLoss(semaphore_0.Wait(1ull, InfiniteFuture()))); // Future flushes/waits/etc should also fail. EXPECT_TRUE(IsDataLoss(command_queue->WaitIdle())); @@ -193,7 +200,7 @@ // Fail. EXPECT_CALL(*mock_target_queue, Submit(_)) .WillOnce([](absl::Span<const SubmissionBatch> batches) { - Sleep(absl::Milliseconds(100)); + Sleep(std::chrono::milliseconds(100)); return DataLossErrorBuilder(IREE_LOC); }); @@ -211,8 +218,8 @@ EXPECT_TRUE(IsDataLoss(command_queue->WaitIdle())); - EXPECT_TRUE(IsDataLoss(semaphore_0.Wait(1ull, absl::InfiniteFuture()))); - EXPECT_TRUE(IsDataLoss(semaphore_1.Wait(1ull, absl::InfiniteFuture()))); + EXPECT_TRUE(IsDataLoss(semaphore_0.Wait(1ull, InfiniteFuture()))); + EXPECT_TRUE(IsDataLoss(semaphore_1.Wait(1ull, InfiniteFuture()))); // Future flushes/waits/etc should also fail. EXPECT_TRUE(IsDataLoss(command_queue->WaitIdle()));
diff --git a/iree/hal/host/serial/serial_scheduling_model.cc b/iree/hal/host/serial/serial_scheduling_model.cc index e343cd6..41cbd9f 100644 --- a/iree/hal/host/serial/serial_scheduling_model.cc +++ b/iree/hal/host/serial/serial_scheduling_model.cc
@@ -58,7 +58,7 @@ return OkStatus(); } - Status WaitIdle(absl::Time deadline) override { + Status WaitIdle(Time deadline_ns) override { // No-op. return OkStatus(); } @@ -110,20 +110,20 @@ } Status SerialSchedulingModel::WaitAllSemaphores( - absl::Span<const SemaphoreValue> semaphores, absl::Time deadline) { + absl::Span<const SemaphoreValue> semaphores, Time deadline_ns) { return CondVarSemaphore::WaitForSemaphores(semaphores, /*wait_all=*/true, - deadline); + deadline_ns); } StatusOr<int> SerialSchedulingModel::WaitAnySemaphore( - absl::Span<const SemaphoreValue> semaphores, absl::Time deadline) { + absl::Span<const SemaphoreValue> semaphores, Time deadline_ns) { return CondVarSemaphore::WaitForSemaphores(semaphores, /*wait_all=*/false, - deadline); + deadline_ns); } -Status SerialSchedulingModel::WaitIdle(absl::Time deadline) { +Status SerialSchedulingModel::WaitIdle(Time deadline_ns) { for (auto& command_queue : command_queues_) { - RETURN_IF_ERROR(command_queue->WaitIdle(deadline)); + RETURN_IF_ERROR(command_queue->WaitIdle(deadline_ns)); } return OkStatus(); }
diff --git a/iree/hal/host/serial/serial_scheduling_model.h b/iree/hal/host/serial/serial_scheduling_model.h index 6aec21f..b065025 100644 --- a/iree/hal/host/serial/serial_scheduling_model.h +++ b/iree/hal/host/serial/serial_scheduling_model.h
@@ -50,10 +50,10 @@ StatusOr<ref_ptr<Semaphore>> CreateSemaphore(uint64_t initial_value) override; Status WaitAllSemaphores(absl::Span<const SemaphoreValue> semaphores, - absl::Time deadline) override; + Time deadline_ns) override; StatusOr<int> WaitAnySemaphore(absl::Span<const SemaphoreValue> semaphores, - absl::Time deadline) override; - Status WaitIdle(absl::Time deadline) override; + Time deadline_ns) override; + Status WaitIdle(Time deadline_ns) override; private: mutable absl::InlinedVector<std::unique_ptr<CommandQueue>, 4> command_queues_;
diff --git a/iree/hal/llvmjit/BUILD b/iree/hal/llvmjit/BUILD index 088bb8b..3ebd609 100644 --- a/iree/hal/llvmjit/BUILD +++ b/iree/hal/llvmjit/BUILD
@@ -64,7 +64,6 @@ srcs = ["llvmjit_executable.cc"], hdrs = ["llvmjit_executable.h"], deps = [ - ":memref_runtime", "//iree/base:status", "//iree/base:tracing", "//iree/hal:buffer", @@ -95,10 +94,3 @@ "//iree/hal:executable_format", ], ) - -cc_library( - name = "memref_runtime", - hdrs = [ - "memref_runtime.h", - ], -)
diff --git a/iree/hal/llvmjit/CMakeLists.txt b/iree/hal/llvmjit/CMakeLists.txt index 8418745..ca40941 100644 --- a/iree/hal/llvmjit/CMakeLists.txt +++ b/iree/hal/llvmjit/CMakeLists.txt
@@ -68,7 +68,6 @@ SRCS "llvmjit_executable.cc" DEPS - ::memref_runtime LLVMAsmParser LLVMCore LLVMOrcJIT @@ -102,11 +101,3 @@ iree::hal::executable_format PUBLIC ) - -iree_cc_library( - NAME - memref_runtime - HDRS - "memref_runtime.h" - PUBLIC -)
diff --git a/iree/hal/llvmjit/llvmjit_executable.cc b/iree/hal/llvmjit/llvmjit_executable.cc index 1596b9e..7d26ccd 100644 --- a/iree/hal/llvmjit/llvmjit_executable.cc +++ b/iree/hal/llvmjit/llvmjit_executable.cc
@@ -21,7 +21,6 @@ #include "iree/base/tracing.h" #include "iree/hal/buffer.h" #include "iree/hal/executable.h" -#include "iree/hal/llvmjit/memref_runtime.h" #include "iree/schemas/llvmir_executable_def_generated.h" #include "llvm/ADT/ArrayRef.h" #include "llvm/ADT/StringRef.h" @@ -82,13 +81,11 @@ make_ref<LLVMJITExecutable>(spec, std::move(ll_jit), allow_aliasing_data); for (const auto func_name : *entry_points) { - auto func_symbol = - executable->ll_jit_->lookup("invoke_" + func_name->str()); + auto func_symbol = executable->ll_jit_->lookup(func_name->str()); if (!func_symbol) { return NotFoundErrorBuilder(IREE_LOC) << "Can't JIT compile function : " << func_name; } - // Map function to its invoke_ symbol. executable->symbols_.push_back(func_symbol.get()); } @@ -111,15 +108,10 @@ struct LLVMJITDispatchState : public HostExecutable::DispatchState { LLVMJITDispatchState() = default; - ~LLVMJITDispatchState() override { - for (int i = 0; i < descriptors.size(); ++i) { - freeUnrankedDescriptor(descriptors[i]); - } - } llvm::JITEvaluatedSymbol symbol; - llvm::SmallVector<UnrankedMemRefType<uint32_t>*, 4> descriptors; llvm::SmallVector<void*, 4> args; + llvm::SmallVector<int64_t, 4> push_constant; }; StatusOr<ref_ptr<HostExecutable::DispatchState>> @@ -142,17 +134,13 @@ MemoryAccessBitfield::kWrite, io_binding.offset, io_binding.length)); auto data = memory.mutable_data(); - auto descriptor = allocUnrankedDescriptor<uint32_t>(data); - dispatch_state->descriptors.push_back(descriptor); - dispatch_state->args.push_back(&descriptor->descriptor); + dispatch_state->args.push_back(data); } } - - auto push_constants_descriptor = allocUnrankedDescriptor<uint32_t>( - const_cast<uint32_t*>(params.push_constants->values.data()), - {static_cast<int64_t>(params.push_constants->values.size())}); - dispatch_state->descriptors.push_back(push_constants_descriptor); - dispatch_state->args.push_back(&push_constants_descriptor->descriptor); + // TODO(ataei): Consider moving this casting to codegen side ?! + for (int i = 0; i < params.push_constants->values.size(); ++i) { + dispatch_state->push_constant.push_back(params.push_constants->values[i]); + } return std::move(dispatch_state); } @@ -162,8 +150,9 @@ IREE_TRACE_SCOPE0("LLVMJITExecutable::DispatchTile"); auto* dispatch_state = static_cast<LLVMJITDispatchState*>(state); - auto func_ptr = (void (*)(void**))dispatch_state->symbol.getAddress(); - func_ptr(dispatch_state->args.data()); + auto func_ptr = + (void (*)(void**, int64_t*))dispatch_state->symbol.getAddress(); + func_ptr(dispatch_state->args.data(), dispatch_state->push_constant.data()); return OkStatus(); }
diff --git a/iree/hal/llvmjit/memref_runtime.h b/iree/hal/llvmjit/memref_runtime.h deleted file mode 100644 index 6b94410..0000000 --- a/iree/hal/llvmjit/memref_runtime.h +++ /dev/null
@@ -1,177 +0,0 @@ -// Copyright 2020 Google LLC -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// https://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// - -#ifndef IREE_HAL_LLVMJIT_LLVMJIT_MEMREF_RUNTIME_H_ -#define IREE_HAL_LLVMJIT_LLVMJIT_MEMREF_RUNTIME_H_ - -#include <assert.h> - -#include <cstdint> -#include <vector> - -namespace iree { -namespace hal { -namespace llvmjit { - -template <int N> -void dropFront(int64_t arr[N], int64_t *res) { - for (unsigned i = 1; i < N; ++i) *(res + i - 1) = arr[i]; -} - -/// StridedMemRef descriptor type with static rank. -template <typename T, int N> -struct StridedMemRefType { - T *basePtr; - T *data; - int64_t offset; - int64_t sizes[N]; - int64_t strides[N]; - // This operator[] is extremely slow and only for sugaring purposes. - StridedMemRefType<T, N - 1> operator[](int64_t idx) { - StridedMemRefType<T, N - 1> res; - res.basePtr = basePtr; - res.data = data; - res.offset = offset + idx * strides[0]; - dropFront<N>(sizes, res.sizes); - dropFront<N>(strides, res.strides); - return res; - } -}; - -/// StridedMemRef descriptor type specialized for rank 1. -template <typename T> -struct StridedMemRefType<T, 1> { - T *basePtr; - T *data; - int64_t offset; - int64_t sizes[1]; - int64_t strides[1]; - T &operator[](int64_t idx) { return *(data + offset + idx * strides[0]); } -}; - -/// StridedMemRef descriptor type specialized for rank 0. -template <typename T> -struct StridedMemRefType<T, 0> { - T *basePtr; - T *data; - int64_t offset; -}; - -// Unranked MemRef -template <typename T> -struct UnrankedMemRefType { - int64_t rank; - void *descriptor; -}; - -// Given a shape with sizes greater than 0 along all dimensions, -// returns the distance, in number of elements, between a slice in a dimension -// and the next slice in the same dimension. -// e.g. shape[3, 4, 5] -> strides[20, 5, 1] -inline std::vector<int64_t> makeStrides(const std::vector<int64_t> &shape) { - std::vector<int64_t> tmp; - if (shape.empty()) return tmp; - tmp.reserve(shape.size()); - int64_t running = 1; - for (auto rit = shape.rbegin(), reit = shape.rend(); rit != reit; ++rit) { - assert(*rit > 0 && - "size must be greater than 0 along all dimensions of shape"); - tmp.push_back(running); - running *= *rit; - } - return std::vector<int64_t>(tmp.rbegin(), tmp.rend()); -} - -// Mallocs a StridedMemRefDescriptor<T, N>* that matches the MLIR ABI. -// This is an implementation detail that is kept in sync with MLIR codegen -// conventions. -template <typename T, int N> -StridedMemRefType<T, N> *makeStridedMemRefDescriptor( - void *ptr, const std::vector<int64_t> &shape) { - StridedMemRefType<T, N> *descriptor = static_cast<StridedMemRefType<T, N> *>( - malloc(sizeof(StridedMemRefType<T, N>))); - descriptor->basePtr = static_cast<T *>(ptr); - descriptor->data = static_cast<T *>(ptr); - descriptor->offset = 0; - std::copy(shape.begin(), shape.end(), descriptor->sizes); - auto strides = makeStrides(shape); - std::copy(strides.begin(), strides.end(), descriptor->strides); - return descriptor; -} - -// Mallocs a StridedMemRefDescriptor<T, 0>* (i.e. a pointer to scalar) that -// matches the MLIR ABI. This is an implementation detail that is kept in sync -// with MLIR codegen conventions. -template <typename T> -StridedMemRefType<T, 0> *makeStridedMemRefDescriptor( - void *ptr, const std::vector<int64_t> &shape) { - StridedMemRefType<T, 0> *descriptor = static_cast<StridedMemRefType<T, 0> *>( - malloc(sizeof(StridedMemRefType<T, 0>))); - descriptor->basePtr = static_cast<T *>(ptr); - descriptor->data = static_cast<T *>(ptr); - descriptor->offset = 0; - return descriptor; -} - -// Mallocs an UnrankedMemRefType<T>* that contains a ranked -// StridedMemRefDescriptor<T, Rank>* and matches the MLIR ABI. This is an -// implementation detail that is kept in sync with MLIR codegen conventions. -template <typename T> -UnrankedMemRefType<T> *allocUnrankedDescriptor( - void *data, const std::vector<int64_t> &shape) { - UnrankedMemRefType<T> *res = static_cast<UnrankedMemRefType<T> *>( - malloc(sizeof(UnrankedMemRefType<T>))); - res->rank = shape.size(); - if (res->rank == 0) - res->descriptor = makeStridedMemRefDescriptor<T>(data, shape); - else if (res->rank == 1) - res->descriptor = makeStridedMemRefDescriptor<T, 1>(data, shape); - else if (res->rank == 2) - res->descriptor = makeStridedMemRefDescriptor<T, 2>(data, shape); - else if (res->rank == 3) - res->descriptor = makeStridedMemRefDescriptor<T, 3>(data, shape); - else if (res->rank == 4) - res->descriptor = makeStridedMemRefDescriptor<T, 4>(data, shape); - else if (res->rank == 5) - res->descriptor = makeStridedMemRefDescriptor<T, 5>(data, shape); - else if (res->rank == 6) - res->descriptor = makeStridedMemRefDescriptor<T, 6>(data, shape); - else - assert(false && "Unsupported 6+D memref descriptor"); - return res; -} - -// Shape and strides aren't used in the generated code (yet). -// TODO(ataei): Delete this version once we can pass shapes. -template <typename T> -UnrankedMemRefType<T> *allocUnrankedDescriptor(void *data) { - UnrankedMemRefType<T> *res = static_cast<UnrankedMemRefType<T> *>( - malloc(sizeof(UnrankedMemRefType<T>))); - res->descriptor = makeStridedMemRefDescriptor<T>(data, {}); - return res; -} - -// Frees an UnrankedMemRefType<T>* -template <typename T> -void freeUnrankedDescriptor(UnrankedMemRefType<T> *desc) { - free(desc->descriptor); - free(desc); -} - -} // namespace llvmjit -} // namespace hal -} // namespace iree - -#endif // IREE_HAL_LLVMJIT_LLVMJIT_MEMREF_RUNTIME_H_
diff --git a/iree/hal/semaphore.h b/iree/hal/semaphore.h index 3a8182c..255585c 100644 --- a/iree/hal/semaphore.h +++ b/iree/hal/semaphore.h
@@ -17,8 +17,6 @@ #include <cstdint> -#include "absl/time/clock.h" -#include "absl/time/time.h" #include "iree/base/status.h" #include "iree/base/time.h" #include "iree/hal/resource.h" @@ -85,16 +83,16 @@ virtual void Fail(Status status) = 0; // Blocks the caller until the semaphore reaches or exceedes the specified - // payload value or the |deadline| elapses. + // payload value or the |deadline_ns| elapses. // // Returns success if the wait is successful and the semaphore has met or // exceeded the required payload value. // - // Returns DEADLINE_EXCEEDED if the |deadline| elapses without the semaphore - // reaching the required value. - virtual Status Wait(uint64_t value, absl::Time deadline) = 0; - inline Status Wait(uint64_t value, absl::Duration timeout) { - return Wait(value, RelativeTimeoutToDeadline(timeout)); + // Returns DEADLINE_EXCEEDED if the |deadline_ns| elapses without the + // semaphore reaching the required value. + virtual Status Wait(uint64_t value, Time deadline_ns) = 0; + inline Status Wait(uint64_t value, Duration timeout_ns) { + return Wait(value, RelativeTimeoutToDeadlineNanos(timeout_ns)); } };
diff --git a/iree/hal/testing/mock_command_queue.h b/iree/hal/testing/mock_command_queue.h index 4530753..c281026 100644 --- a/iree/hal/testing/mock_command_queue.h +++ b/iree/hal/testing/mock_command_queue.h
@@ -32,7 +32,7 @@ MOCK_METHOD(Status, Submit, (absl::Span<const SubmissionBatch> batches), (override)); - MOCK_METHOD(Status, WaitIdle, (absl::Time deadline), (override)); + MOCK_METHOD(Status, WaitIdle, (Time deadline_ns), (override)); }; } // namespace testing
diff --git a/iree/hal/vmla/BUILD b/iree/hal/vmla/BUILD index d966b7c..35e081a 100644 --- a/iree/hal/vmla/BUILD +++ b/iree/hal/vmla/BUILD
@@ -140,8 +140,8 @@ "//iree/vm:context", "//iree/vm:instance", "//iree/vm:invocation", + "//iree/vm:list", "//iree/vm:module", - "//iree/vm:variant_list", "@com_google_absl//absl/container:inlined_vector", "@com_google_absl//absl/types:span", ],
diff --git a/iree/hal/vmla/CMakeLists.txt b/iree/hal/vmla/CMakeLists.txt index adb4626..cfa8cca 100644 --- a/iree/hal/vmla/CMakeLists.txt +++ b/iree/hal/vmla/CMakeLists.txt
@@ -150,8 +150,8 @@ iree::vm::context iree::vm::instance iree::vm::invocation + iree::vm::list iree::vm::module - iree::vm::variant_list PUBLIC )
diff --git a/iree/hal/vmla/vmla_executable.cc b/iree/hal/vmla/vmla_executable.cc index 2990219..489a215 100644 --- a/iree/hal/vmla/vmla_executable.cc +++ b/iree/hal/vmla/vmla_executable.cc
@@ -22,8 +22,8 @@ #include "iree/schemas/vmla_executable_def_generated.h" #include "iree/vm/bytecode_module.h" #include "iree/vm/invocation.h" +#include "iree/vm/list.h" #include "iree/vm/module.h" -#include "iree/vm/variant_list.h" namespace iree { namespace hal { @@ -133,8 +133,8 @@ auto dispatch_state = make_ref<VMLADispatchState>(); dispatch_state->function = entry_functions_[params.entry_point]; - dispatch_state->input_list_size = - iree_vm_variant_list_alloc_size(/*interface*/ 1 + /*workgroup_xyz[3]*/ 3); + dispatch_state->input_list_size = iree_vm_list_storage_size( + /*element_type=*/nullptr, /*interface*/ 1 + /*workgroup_xyz[3]*/ 3); auto* interface = &dispatch_state->interface; RETURN_IF_ERROR(interface->SetConstants(params.push_constants->values)); @@ -163,18 +163,20 @@ IREE_TRACE_SCOPE0("VMLAExecutable::DispatchTile"); auto* dispatch_state = static_cast<VMLADispatchState*>(state); - auto* input_list = reinterpret_cast<iree_vm_variant_list_t*>( - alloca(dispatch_state->input_list_size)); - iree_vm_variant_list_init(input_list, - /*interface*/ 1 + /*workgroup_xyz[3]*/ 3); - iree_vm_variant_list_append_ref_retain(input_list, - &dispatch_state->interface_ref); - iree_vm_variant_list_append_value(input_list, - iree_vm_value_make_i32(workgroup_xyz[0])); - iree_vm_variant_list_append_value(input_list, - iree_vm_value_make_i32(workgroup_xyz[1])); - iree_vm_variant_list_append_value(input_list, - iree_vm_value_make_i32(workgroup_xyz[2])); + auto* input_list_storage = alloca(dispatch_state->input_list_size); + iree_vm_list_t* input_list = nullptr; + RETURN_IF_ERROR( + FromApiStatus(iree_vm_list_initialize( + iree_make_byte_span(input_list_storage, + dispatch_state->input_list_size), + /*element_type=*/nullptr, + /*interface*/ 1 + /*workgroup_xyz[3]*/ 3, &input_list), + IREE_LOC)); + iree_vm_list_push_ref_retain(input_list, &dispatch_state->interface_ref); + for (int i = 0; i < workgroup_xyz.size(); ++i) { + iree_vm_value_t value = iree_vm_value_make_i32(workgroup_xyz[i]); + iree_vm_list_push_value(input_list, &value); + } auto status = FromApiStatus(iree_vm_invoke(context(), dispatch_state->function, @@ -182,7 +184,7 @@ /*outputs=*/nullptr, IREE_ALLOCATOR_SYSTEM), IREE_LOC); - iree_vm_variant_list_free(input_list); + iree_vm_list_deinitialize(input_list); return std::move(status); }
diff --git a/iree/hal/vmla/vmla_executable.h b/iree/hal/vmla/vmla_executable.h index d74b11a..0e79b36 100644 --- a/iree/hal/vmla/vmla_executable.h +++ b/iree/hal/vmla/vmla_executable.h
@@ -25,7 +25,6 @@ #include "iree/vm/context.h" #include "iree/vm/instance.h" #include "iree/vm/module.h" -#include "iree/vm/variant_list.h" namespace iree { namespace hal {
diff --git a/iree/hal/vulkan/BUILD b/iree/hal/vulkan/BUILD index b93167a..ca73211 100644 --- a/iree/hal/vulkan/BUILD +++ b/iree/hal/vulkan/BUILD
@@ -138,6 +138,7 @@ ":handle_util", ":native_timeline_semaphore", ":status_util", + "//iree/base:api", "//iree/base:arena", "//iree/base:memory", "//iree/base:source_location", @@ -146,7 +147,6 @@ "//iree/hal:command_queue", "@com_google_absl//absl/base:core_headers", "@com_google_absl//absl/synchronization", - "@com_google_absl//absl/time", "@iree_vulkan_headers//:vulkan_headers_no_prototypes", ], ) @@ -194,15 +194,16 @@ ":handle_util", ":status_util", ":timepoint_util", + "//iree/base:api", "//iree/base:intrusive_list", "//iree/base:ref_ptr", "//iree/base:status", + "//iree/base:time", "//iree/base:tracing", "//iree/hal:semaphore", "@com_google_absl//absl/base:core_headers", "@com_google_absl//absl/container:inlined_vector", "@com_google_absl//absl/synchronization", - "@com_google_absl//absl/time", "@com_google_absl//absl/utility", "@iree_vulkan_headers//:vulkan_headers_no_prototypes", ], @@ -358,6 +359,7 @@ ":handle_util", ":status_util", ":timepoint_util", + "//iree/base:api", "//iree/base:intrusive_list", "//iree/base:memory", "//iree/base:ref_ptr", @@ -370,7 +372,6 @@ "@com_google_absl//absl/base:core_headers", "@com_google_absl//absl/container:inlined_vector", "@com_google_absl//absl/synchronization", - "@com_google_absl//absl/time", "@com_google_absl//absl/types:span", ], ) @@ -396,10 +397,10 @@ "//iree/base:intrusive_list", "//iree/base:ref_ptr", "//iree/base:status", + "//iree/base:time", "//iree/base:tracing", "@com_google_absl//absl/base:core_headers", "@com_google_absl//absl/synchronization", - "@com_google_absl//absl/time", "@com_google_absl//absl/utility", "@iree_vulkan_headers//:vulkan_headers_no_prototypes", ], @@ -468,6 +469,7 @@ "//iree/base:math", "//iree/base:memory", "//iree/base:status", + "//iree/base:time", "//iree/base:tracing", "//iree/hal:allocator", "//iree/hal:command_buffer_validation",
diff --git a/iree/hal/vulkan/CMakeLists.txt b/iree/hal/vulkan/CMakeLists.txt index e69846d..78e81f7 100644 --- a/iree/hal/vulkan/CMakeLists.txt +++ b/iree/hal/vulkan/CMakeLists.txt
@@ -141,10 +141,10 @@ DEPS absl::core_headers absl::synchronization - absl::time iree::base::arena iree::base::memory iree::base::status + iree::base::time iree::base::tracing iree::hal::command_queue iree::hal::vulkan::direct_command_buffer @@ -212,9 +212,9 @@ ::timepoint_util absl::inlined_vector absl::synchronization - absl::time iree::base::intrusive_list iree::base::status + iree::base::time iree::base::tracing iree::hal::semaphore Vulkan::Headers @@ -454,6 +454,7 @@ iree::base::intrusive_list iree::base::ref_ptr iree::base::status + iree::base::time iree::base::tracing Vulkan::Headers PUBLIC @@ -526,6 +527,7 @@ iree::base::math iree::base::memory iree::base::status + iree::base::time iree::base::tracing iree::hal::allocator iree::hal::command_buffer_validation
diff --git a/iree/hal/vulkan/direct_command_queue.cc b/iree/hal/vulkan/direct_command_queue.cc index 11a33cc..84f3876 100644 --- a/iree/hal/vulkan/direct_command_queue.cc +++ b/iree/hal/vulkan/direct_command_queue.cc
@@ -16,8 +16,6 @@ #include <cstdint> -#include "absl/time/clock.h" -#include "absl/time/time.h" #include "iree/base/memory.h" #include "iree/base/source_location.h" #include "iree/base/status.h" @@ -134,8 +132,8 @@ return OkStatus(); } -Status DirectCommandQueue::WaitIdle(absl::Time deadline) { - if (deadline == absl::InfiniteFuture()) { +Status DirectCommandQueue::WaitIdle(Time deadline_ns) { + if (deadline_ns == InfiniteFuture()) { // Fast path for using vkQueueWaitIdle, which is usually cheaper (as it // requires fewer calls into the driver). IREE_TRACE_SCOPE0("DirectCommandQueue::WaitIdle#vkQueueWaitIdle"); @@ -160,21 +158,21 @@ logical_device_->allocator()); }); - uint64_t timeout; - if (deadline == absl::InfinitePast()) { + uint64_t timeout_ns; + if (deadline_ns == InfinitePast()) { // Do not wait. - timeout = 0; - } else if (deadline == absl::InfiniteFuture()) { + timeout_ns = 0; + } else if (deadline_ns == InfiniteFuture()) { // Wait forever. - timeout = UINT64_MAX; + timeout_ns = UINT64_MAX; } else { // Convert to relative time in nanoseconds. // The implementation may not wait with this granularity (like, by 10000x). - absl::Time now = absl::Now(); - if (deadline < now) { + Time now_ns = Now(); + if (deadline_ns < now_ns) { return DeadlineExceededErrorBuilder(IREE_LOC) << "Deadline in the past"; } - timeout = static_cast<uint64_t>(absl::ToInt64Nanoseconds(deadline - now)); + timeout_ns = static_cast<uint64_t>(deadline_ns - now_ns); } { @@ -183,7 +181,7 @@ } VkResult result = - syms()->vkWaitForFences(*logical_device_, 1, &fence, VK_TRUE, timeout); + syms()->vkWaitForFences(*logical_device_, 1, &fence, VK_TRUE, timeout_ns); switch (result) { case VK_SUCCESS: return OkStatus();
diff --git a/iree/hal/vulkan/direct_command_queue.h b/iree/hal/vulkan/direct_command_queue.h index c080398..1921bfa 100644 --- a/iree/hal/vulkan/direct_command_queue.h +++ b/iree/hal/vulkan/direct_command_queue.h
@@ -22,9 +22,9 @@ #include "absl/base/thread_annotations.h" #include "absl/synchronization/mutex.h" -#include "absl/time/time.h" #include "iree/base/arena.h" #include "iree/base/status.h" +#include "iree/base/time.h" #include "iree/hal/command_queue.h" #include "iree/hal/vulkan/dynamic_symbols.h" #include "iree/hal/vulkan/handle_util.h" @@ -48,7 +48,7 @@ Status Submit(absl::Span<const SubmissionBatch> batches) override; - Status WaitIdle(absl::Time deadline) override; + Status WaitIdle(Time deadline_ns) override; private: Status TranslateBatchInfo(const SubmissionBatch& batch,
diff --git a/iree/hal/vulkan/emulated_timeline_semaphore.cc b/iree/hal/vulkan/emulated_timeline_semaphore.cc index 9656310..a3e2c1e 100644 --- a/iree/hal/vulkan/emulated_timeline_semaphore.cc +++ b/iree/hal/vulkan/emulated_timeline_semaphore.cc
@@ -16,8 +16,8 @@ #include "absl/container/inlined_vector.h" #include "absl/synchronization/mutex.h" -#include "absl/time/time.h" #include "absl/utility/utility.h" +#include "iree/base/time.h" #include "iree/base/tracing.h" #include "iree/hal/vulkan/dynamic_symbols.h" #include "iree/hal/vulkan/status_util.h" @@ -82,7 +82,7 @@ return OkStatus(); } -Status EmulatedTimelineSemaphore::Wait(uint64_t value, absl::Time deadline) { +Status EmulatedTimelineSemaphore::Wait(uint64_t value, Time deadline_ns) { IREE_TRACE_SCOPE0("EmulatedTimelineSemaphore::Wait"); VkFence fence = VK_NULL_HANDLE; @@ -112,26 +112,18 @@ break; } // TODO(antiagainst): figure out a better way instead of the busy loop here. - } while (absl::Now() < deadline); + } while (Now() < deadline_ns); if (fence == VK_NULL_HANDLE) { return DeadlineExceededErrorBuilder(IREE_LOC) << "Deadline reached when waiting timeline semaphore"; } - uint64_t timeout_nanos; - if (deadline == absl::InfiniteFuture()) { - timeout_nanos = UINT64_MAX; - } else if (deadline == absl::InfinitePast()) { - timeout_nanos = 0; - } else { - auto relative_nanos = absl::ToInt64Nanoseconds(deadline - absl::Now()); - timeout_nanos = relative_nanos < 0 ? 0 : relative_nanos; - } - + uint64_t timeout_ns = + static_cast<uint64_t>(DeadlineToRelativeTimeoutNanos(deadline_ns)); VK_RETURN_IF_ERROR(logical_device_->syms()->vkWaitForFences( *logical_device_, /*fenceCount=*/1, &fence, /*waitAll=*/true, - timeout_nanos)); + timeout_ns)); RETURN_IF_ERROR(TryToAdvanceTimeline(value).status()); return OkStatus();
diff --git a/iree/hal/vulkan/emulated_timeline_semaphore.h b/iree/hal/vulkan/emulated_timeline_semaphore.h index cc13a09..d4d61c8 100644 --- a/iree/hal/vulkan/emulated_timeline_semaphore.h +++ b/iree/hal/vulkan/emulated_timeline_semaphore.h
@@ -144,7 +144,7 @@ Status Signal(uint64_t value) override; - Status Wait(uint64_t value, absl::Time deadline) override; + Status Wait(uint64_t value, Time deadline_ns) override; void Fail(Status status) override;
diff --git a/iree/hal/vulkan/native_timeline_semaphore.cc b/iree/hal/vulkan/native_timeline_semaphore.cc index c0898b4..2a9ba8b 100644 --- a/iree/hal/vulkan/native_timeline_semaphore.cc +++ b/iree/hal/vulkan/native_timeline_semaphore.cc
@@ -97,7 +97,7 @@ logical_device_->syms()->vkSignalSemaphore(*logical_device_, &signal_info); } -Status NativeTimelineSemaphore::Wait(uint64_t value, absl::Time deadline) { +Status NativeTimelineSemaphore::Wait(uint64_t value, Time deadline_ns) { IREE_TRACE_SCOPE0("NativeTimelineSemaphore::Wait"); VkSemaphoreWaitInfo wait_info; @@ -108,14 +108,15 @@ wait_info.pSemaphores = &handle_; wait_info.pValues = &value; - uint64_t timeout_nanos; - if (deadline == absl::InfiniteFuture()) { - timeout_nanos = UINT64_MAX; - } else if (deadline == absl::InfinitePast()) { - timeout_nanos = 0; + uint64_t timeout_ns; + if (deadline_ns == InfiniteFuture()) { + timeout_ns = UINT64_MAX; + } else if (deadline_ns == InfinitePast()) { + timeout_ns = 0; } else { - auto relative_nanos = absl::ToInt64Nanoseconds(deadline - absl::Now()); - timeout_nanos = relative_nanos < 0 ? 0 : relative_nanos; + Duration relative_ns = deadline_ns - Now(); + timeout_ns = static_cast<int64_t>( + relative_ns < ZeroDuration() ? ZeroDuration() : relative_ns); } // NOTE: this may fail with a timeout (VK_TIMEOUT) or in the case of a @@ -126,7 +127,7 @@ return UnknownErrorBuilder(IREE_LOC) << "vkWaitSemaphores not defined"; } VkResult result = logical_device_->syms()->vkWaitSemaphores( - *logical_device_, &wait_info, timeout_nanos); + *logical_device_, &wait_info, timeout_ns); if (result == VK_ERROR_DEVICE_LOST) { // Nothing we do now matters. return VkResultToStatus(result);
diff --git a/iree/hal/vulkan/native_timeline_semaphore.h b/iree/hal/vulkan/native_timeline_semaphore.h index 3f44351..a5d3a93 100644 --- a/iree/hal/vulkan/native_timeline_semaphore.h +++ b/iree/hal/vulkan/native_timeline_semaphore.h
@@ -44,7 +44,7 @@ Status Signal(uint64_t value) override; void Fail(Status status) override; - Status Wait(uint64_t value, absl::Time deadline) override; + Status Wait(uint64_t value, Time deadline_ns) override; private: ref_ptr<VkDeviceHandle> logical_device_;
diff --git a/iree/hal/vulkan/serializing_command_queue.cc b/iree/hal/vulkan/serializing_command_queue.cc index 563f4c7..3a49ba2 100644 --- a/iree/hal/vulkan/serializing_command_queue.cc +++ b/iree/hal/vulkan/serializing_command_queue.cc
@@ -16,8 +16,8 @@ #include <memory> -#include "absl/time/clock.h" #include "absl/types/span.h" +#include "iree/base/api.h" #include "iree/base/memory.h" #include "iree/base/source_location.h" #include "iree/base/tracing.h" @@ -266,10 +266,10 @@ return true; } -Status SerializingCommandQueue::WaitIdle(absl::Time deadline) { +Status SerializingCommandQueue::WaitIdle(Time deadline_ns) { absl::MutexLock lock(&mutex_); - if (deadline == absl::InfiniteFuture()) { + if (deadline_ns == InfiniteFuture()) { IREE_TRACE_SCOPE0("SerializingCommandQueue::WaitIdle#vkQueueWaitIdle"); // Fast path for using vkQueueWaitIdle, which is usually cheaper (as it // requires fewer calls into the driver). @@ -296,21 +296,21 @@ do { RETURN_IF_ERROR(ProcessDeferredSubmissions().status()); - uint64_t timeout_nanos; - if (deadline == absl::InfinitePast()) { - // Do not wait. - timeout_nanos = 0; + uint64_t timeout_ns; + if (deadline_ns == InfiniteFuture()) { + timeout_ns = UINT64_MAX; + } else if (deadline_ns == InfinitePast()) { + timeout_ns = 0; } else { // Convert to relative time in nanoseconds. // The implementation may not wait with this granularity (like, by // 10000x). - absl::Time now = absl::Now(); - if (deadline < now) { + Duration relative_ns = deadline_ns - Now(); + if (relative_ns < ZeroDuration()) { return DeadlineExceededErrorBuilder(IREE_LOC) << "Deadline exceeded waiting for idle"; } - timeout_nanos = - static_cast<uint64_t>(absl::ToInt64Nanoseconds(deadline - now)); + timeout_ns = static_cast<uint64_t>(relative_ns); } if (pending_fences_.empty()) continue; @@ -321,7 +321,7 @@ VkResult result = syms()->vkWaitForFences(*logical_device_, fences.size(), fences.data(), - /*waitAll=*/VK_TRUE, timeout_nanos); + /*waitAll=*/VK_TRUE, timeout_ns); switch (result) { case VK_SUCCESS:
diff --git a/iree/hal/vulkan/serializing_command_queue.h b/iree/hal/vulkan/serializing_command_queue.h index e38643b..92a811d 100644 --- a/iree/hal/vulkan/serializing_command_queue.h +++ b/iree/hal/vulkan/serializing_command_queue.h
@@ -23,10 +23,10 @@ #include "absl/base/thread_annotations.h" #include "absl/container/inlined_vector.h" #include "absl/synchronization/mutex.h" -#include "absl/time/time.h" #include "iree/base/intrusive_list.h" #include "iree/base/ref_ptr.h" #include "iree/base/status.h" +#include "iree/base/time.h" #include "iree/hal/command_buffer.h" #include "iree/hal/command_queue.h" #include "iree/hal/vulkan/dynamic_symbols.h" @@ -63,7 +63,7 @@ Status Submit(absl::Span<const SubmissionBatch> batches) override; - Status WaitIdle(absl::Time deadline) override; + Status WaitIdle(Time deadline_ns) override; // Releases all deferred submissions ready to submit to the GPU. Status AdvanceQueueSubmission();
diff --git a/iree/hal/vulkan/timepoint_util.cc b/iree/hal/vulkan/timepoint_util.cc index 98e9f96..7664f3b 100644 --- a/iree/hal/vulkan/timepoint_util.cc +++ b/iree/hal/vulkan/timepoint_util.cc
@@ -17,8 +17,8 @@ #include <memory> #include "absl/synchronization/mutex.h" -#include "absl/time/time.h" #include "absl/utility/utility.h" +#include "iree/base/time.h" #include "iree/base/tracing.h" #include "iree/hal/vulkan/dynamic_symbols.h" #include "iree/hal/vulkan/status_util.h"
diff --git a/iree/hal/vulkan/vulkan_device.cc b/iree/hal/vulkan/vulkan_device.cc index 55a4a4e..3209d5a 100644 --- a/iree/hal/vulkan/vulkan_device.cc +++ b/iree/hal/vulkan/vulkan_device.cc
@@ -24,6 +24,7 @@ #include "absl/synchronization/mutex.h" #include "iree/base/math.h" #include "iree/base/status.h" +#include "iree/base/time.h" #include "iree/base/tracing.h" #include "iree/hal/command_buffer_validation.h" #include "iree/hal/command_queue.h" @@ -736,20 +737,20 @@ } Status VulkanDevice::WaitAllSemaphores( - absl::Span<const SemaphoreValue> semaphores, absl::Time deadline) { + absl::Span<const SemaphoreValue> semaphores, Time deadline_ns) { IREE_TRACE_SCOPE0("VulkanDevice::WaitAllSemaphores"); - return WaitSemaphores(semaphores, deadline, /*wait_flags=*/0); + return WaitSemaphores(semaphores, deadline_ns, /*wait_flags=*/0); } StatusOr<int> VulkanDevice::WaitAnySemaphore( - absl::Span<const SemaphoreValue> semaphores, absl::Time deadline) { + absl::Span<const SemaphoreValue> semaphores, Time deadline_ns) { IREE_TRACE_SCOPE0("VulkanDevice::WaitAnySemaphore"); - return WaitSemaphores(semaphores, deadline, + return WaitSemaphores(semaphores, deadline_ns, /*wait_flags=*/VK_SEMAPHORE_WAIT_ANY_BIT); } Status VulkanDevice::WaitSemaphores(absl::Span<const SemaphoreValue> semaphores, - absl::Time deadline, + Time deadline_ns, VkSemaphoreWaitFlags wait_flags) { IREE_TRACE_SCOPE0("VulkanDevice::WaitSemaphores"); @@ -762,7 +763,7 @@ for (int i = 0; i < semaphores.size(); ++i) { auto* semaphore = static_cast<EmulatedTimelineSemaphore*>(semaphores[i].semaphore); - RETURN_IF_ERROR(semaphore->Wait(semaphores[i].value, deadline)); + RETURN_IF_ERROR(semaphore->Wait(semaphores[i].value, deadline_ns)); if (wait_flags & VK_SEMAPHORE_WAIT_ANY_BIT) return OkStatus(); } @@ -786,22 +787,14 @@ wait_info.pSemaphores = semaphore_handles.data(); wait_info.pValues = semaphore_values.data(); - uint64_t timeout_nanos; - if (deadline == absl::InfiniteFuture()) { - timeout_nanos = UINT64_MAX; - } else if (deadline == absl::InfinitePast()) { - timeout_nanos = 0; - } else { - auto relative_nanos = absl::ToInt64Nanoseconds(deadline - absl::Now()); - timeout_nanos = relative_nanos < 0 ? 0 : relative_nanos; - } - // NOTE: this may fail with a timeout (VK_TIMEOUT) or in the case of a // device loss event may return either VK_SUCCESS *or* VK_ERROR_DEVICE_LOST. // We may want to explicitly query for device loss after a successful wait // to ensure we consistently return errors. + uint64_t timeout_ns = + static_cast<uint64_t>(DeadlineToRelativeTimeoutNanos(deadline_ns)); VkResult result = - syms()->vkWaitSemaphores(*logical_device_, &wait_info, timeout_nanos); + syms()->vkWaitSemaphores(*logical_device_, &wait_info, timeout_ns); if (result == VK_ERROR_DEVICE_LOST) { // Nothing we do now matters. return VkResultToStatus(result); @@ -813,8 +806,8 @@ return OkStatus(); } -Status VulkanDevice::WaitIdle(absl::Time deadline) { - if (deadline == absl::InfiniteFuture()) { +Status VulkanDevice::WaitIdle(Time deadline_ns) { + if (deadline_ns == InfiniteFuture()) { // Fast path for using vkDeviceWaitIdle, which is usually cheaper (as it // requires fewer calls into the driver). IREE_TRACE_SCOPE0("VulkanDevice::WaitIdle#vkDeviceWaitIdle"); @@ -824,7 +817,7 @@ IREE_TRACE_SCOPE0("VulkanDevice::WaitIdle#Semaphores"); for (auto& command_queue : command_queues_) { - RETURN_IF_ERROR(command_queue->WaitIdle(deadline)); + RETURN_IF_ERROR(command_queue->WaitIdle(deadline_ns)); } return OkStatus(); }
diff --git a/iree/hal/vulkan/vulkan_device.h b/iree/hal/vulkan/vulkan_device.h index 4032bbf..273d407 100644 --- a/iree/hal/vulkan/vulkan_device.h +++ b/iree/hal/vulkan/vulkan_device.h
@@ -106,11 +106,11 @@ StatusOr<ref_ptr<Semaphore>> CreateSemaphore(uint64_t initial_value) override; Status WaitAllSemaphores(absl::Span<const SemaphoreValue> semaphores, - absl::Time deadline) override; + Time deadline_ns) override; StatusOr<int> WaitAnySemaphore(absl::Span<const SemaphoreValue> semaphores, - absl::Time deadline) override; + Time deadline_ns) override; - Status WaitIdle(absl::Time deadline) override; + Status WaitIdle(Time deadline_ns) override; private: VulkanDevice( @@ -125,7 +125,7 @@ DebugCaptureManager* debug_capture_manager); Status WaitSemaphores(absl::Span<const SemaphoreValue> semaphores, - absl::Time deadline, VkSemaphoreWaitFlags wait_flags); + Time deadline_ns, VkSemaphoreWaitFlags wait_flags); bool emulating_timeline_semaphores() const { return semaphore_pool_ != nullptr;
diff --git a/iree/modules/check/BUILD b/iree/modules/check/BUILD index 863b189..f1e5cbf 100644 --- a/iree/modules/check/BUILD +++ b/iree/modules/check/BUILD
@@ -41,7 +41,6 @@ "//iree/testing:gtest_main", "//iree/vm", "//iree/vm:bytecode_module", - "//iree/vm:ref", "//iree/vm:ref_cc", "@com_google_absl//absl/base:core_headers", "@com_google_absl//absl/strings",
diff --git a/iree/modules/check/CMakeLists.txt b/iree/modules/check/CMakeLists.txt index 804b844..0a016be 100644 --- a/iree/modules/check/CMakeLists.txt +++ b/iree/modules/check/CMakeLists.txt
@@ -34,7 +34,6 @@ iree::testing::gtest_main iree::vm iree::vm::bytecode_module - iree::vm::ref iree::vm::ref_cc )
diff --git a/iree/modules/check/check_test.cc b/iree/modules/check/check_test.cc index c1cb470..ed84a71 100644 --- a/iree/modules/check/check_test.cc +++ b/iree/modules/check/check_test.cc
@@ -27,7 +27,6 @@ #include "iree/testing/gtest.h" #include "iree/vm/api.h" #include "iree/vm/bytecode_module.h" -#include "iree/vm/ref.h" #include "iree/vm/ref_cc.h" namespace iree { @@ -71,8 +70,8 @@ } void TearDown() override { + inputs_.reset(); iree_vm_context_release(context_); - if (inputs_) iree_vm_variant_list_free(inputs_); } void CreateInt32BufferView(absl::Span<const int32_t> contents, @@ -173,7 +172,7 @@ // TODO(#2075): don't directly invoke native functions like this. return FromApiStatus( iree_vm_invoke(context_, function, - /*policy=*/nullptr, inputs_, + /*policy=*/nullptr, inputs_.get(), /*outputs=*/nullptr, IREE_ALLOCATOR_SYSTEM), IREE_LOC); } @@ -181,12 +180,12 @@ Status Invoke(absl::string_view function_name, std::vector<iree_vm_value> args) { RETURN_IF_ERROR( - FromApiStatus(iree_vm_variant_list_alloc( - args.size(), IREE_ALLOCATOR_SYSTEM, &inputs_), + FromApiStatus(iree_vm_list_create(/*element_type=*/nullptr, args.size(), + IREE_ALLOCATOR_SYSTEM, &inputs_), IREE_LOC)); for (iree_vm_value& arg : args) { RETURN_IF_ERROR(FromApiStatus( - iree_vm_variant_list_append_value(inputs_, arg), IREE_LOC)); + iree_vm_list_push_value(inputs_.get(), &arg), IREE_LOC)); } return Invoke(function_name); } @@ -194,13 +193,13 @@ Status Invoke(absl::string_view function_name, std::vector<vm::ref<iree_hal_buffer_view_t>> args) { RETURN_IF_ERROR( - FromApiStatus(iree_vm_variant_list_alloc( - args.size(), IREE_ALLOCATOR_SYSTEM, &inputs_), + FromApiStatus(iree_vm_list_create(/*element_type=*/nullptr, args.size(), + IREE_ALLOCATOR_SYSTEM, &inputs_), IREE_LOC)); for (auto& arg : args) { iree_vm_ref_t arg_ref = iree_hal_buffer_view_move_ref(arg.get()); RETURN_IF_ERROR(FromApiStatus( - iree_vm_variant_list_append_ref_retain(inputs_, &arg_ref), IREE_LOC)); + iree_vm_list_push_ref_retain(inputs_.get(), &arg_ref), IREE_LOC)); } return Invoke(function_name); } @@ -212,7 +211,7 @@ static iree_vm_module_t* hal_module_; iree_vm_context_t* context_ = nullptr; - iree_vm_variant_list_t* inputs_ = nullptr; + vm::ref<iree_vm_list_t> inputs_; iree_hal_allocator_t* allocator_ = nullptr; }; iree_hal_device_t* CheckTest::device_ = nullptr;
diff --git a/iree/modules/hal/hal_module.cc b/iree/modules/hal/hal_module.cc index 2b1d6db..e51938f 100644 --- a/iree/modules/hal/hal_module.cc +++ b/iree/modules/hal/hal_module.cc
@@ -457,6 +457,30 @@ return BufferViewDimsN<4>(std::move(buffer_view)); } + Status BufferViewTrace( + absl::Span<const vm::ref<iree_hal_buffer_view_t>> buffer_views) { + // TODO(hanchung): Have better information for each dump, eg, having StrAttr + // for each trace event so we can map the dump to dispatch functions easier. + fprintf(stderr, "=== DEBUG DUMP ===\n"); + for (auto& view : buffer_views) { + std::string result_str(4096, '\0'); + iree_status_t status; + do { + iree_host_size_t actual_length = 0; + status = iree_hal_buffer_view_format( + view.get(), /*max_element_count=*/1024, result_str.size() + 1, + &result_str[0], &actual_length); + result_str.resize(actual_length); + } while (iree_status_is_out_of_range(status)); + if (!iree_status_is_ok(status)) { + return FromApiStatus(status, IREE_LOC); + } + fprintf(stderr, "%s\n", result_str.c_str()); + } + fprintf(stderr, "\n"); + return OkStatus(); + } + //===--------------------------------------------------------------------===// // iree::hal::CommandBuffer //===--------------------------------------------------------------------===// @@ -876,6 +900,8 @@ &HALModuleState::BufferViewDims3), vm::MakeNativeFunction("buffer_view.dims.4", &HALModuleState::BufferViewDims4), + vm::MakeNativeFunction("buffer_view.trace", + &HALModuleState::BufferViewTrace), vm::MakeNativeFunction("command_buffer.create", &HALModuleState::CommandBufferCreate),
diff --git a/iree/modules/strings/BUILD b/iree/modules/strings/BUILD index 4d46042..d4e7053 100644 --- a/iree/modules/strings/BUILD +++ b/iree/modules/strings/BUILD
@@ -27,7 +27,6 @@ "//iree/vm:module", "//iree/vm:module_abi_cc", "//iree/vm:ref", - "//iree/vm:ref_cc", "//iree/vm:stack", "@com_google_absl//absl/container:inlined_vector", "@com_google_absl//absl/strings", @@ -48,13 +47,10 @@ "//iree/hal/vmla:vmla_driver_module", "//iree/modules/hal", "//iree/testing:gtest_main", + "//iree/vm", "//iree/vm:bytecode_module", - "//iree/vm:context", - "//iree/vm:instance", - "//iree/vm:module", - "//iree/vm:ref", + "//iree/vm:module_abi_cc", "//iree/vm:ref_cc", - "//iree/vm:stack", "@com_google_absl//absl/container:inlined_vector", "@com_google_absl//absl/strings", "@com_google_benchmark//:benchmark",
diff --git a/iree/modules/strings/CMakeLists.txt b/iree/modules/strings/CMakeLists.txt index 0c898bb..6c6ac35 100644 --- a/iree/modules/strings/CMakeLists.txt +++ b/iree/modules/strings/CMakeLists.txt
@@ -38,7 +38,6 @@ iree::vm::module iree::vm::module_abi_cc iree::vm::ref - iree::vm::ref_cc iree::vm::stack PUBLIC ) @@ -60,13 +59,10 @@ iree::hal::vmla::vmla_driver_module iree::modules::hal iree::testing::gtest_main + iree::vm iree::vm::bytecode_module - iree::vm::context - iree::vm::instance - iree::vm::module - iree::vm::ref + iree::vm::module_abi_cc iree::vm::ref_cc - iree::vm::stack ) iree_bytecode_module(
diff --git a/iree/modules/strings/strings_module.cc b/iree/modules/strings/strings_module.cc index a52bac0..abd07ee 100644 --- a/iree/modules/strings/strings_module.cc +++ b/iree/modules/strings/strings_module.cc
@@ -28,10 +28,7 @@ #include "iree/modules/strings/api.h" #include "iree/modules/strings/api_detail.h" #include "iree/vm/bytecode_module.h" -#include "iree/vm/module.h" #include "iree/vm/module_abi_cc.h" -#include "iree/vm/ref.h" -#include "iree/vm/stack.h" static iree_vm_ref_type_descriptor_t strings_string_descriptor = {0}; static iree_vm_ref_type_descriptor_t strings_string_tensor_descriptor = {0};
diff --git a/iree/modules/strings/strings_module_test.cc b/iree/modules/strings/strings_module_test.cc index b3a3122..d5ee15d 100644 --- a/iree/modules/strings/strings_module_test.cc +++ b/iree/modules/strings/strings_module_test.cc
@@ -26,13 +26,9 @@ #include "iree/modules/strings/api_detail.h" #include "iree/modules/strings/strings_module_test_module.h" #include "iree/testing/gtest.h" +#include "iree/vm/api.h" #include "iree/vm/bytecode_module.h" -#include "iree/vm/context.h" -#include "iree/vm/instance.h" -#include "iree/vm/module.h" -#include "iree/vm/ref.h" #include "iree/vm/ref_cc.h" -#include "iree/vm/stack.h" using testing::internal::CaptureStdout; using testing::internal::GetCapturedStdout; @@ -138,37 +134,33 @@ shape.data(), shape.size(), &input_string_tensor)); // Construct the input list for execution. - iree_vm_variant_list_t* inputs = nullptr; - IREE_ASSERT_OK( - iree_vm_variant_list_alloc(1, IREE_ALLOCATOR_SYSTEM, &inputs)); + vm::ref<iree_vm_list_t> inputs; + IREE_ASSERT_OK(iree_vm_list_create(/*element_type=*/nullptr, 1, + IREE_ALLOCATOR_SYSTEM, &inputs)); // Add the string tensor to the input list. - iree_vm_ref_t input_string_tensor_ref = - strings_string_tensor_move_ref(input_string_tensor.get()); - IREE_ASSERT_OK(iree_vm_variant_list_append_ref_retain( - inputs, &input_string_tensor_ref)); + IREE_ASSERT_OK( + iree_vm_list_push_ref_retain(inputs.get(), input_string_tensor)); // Construct the output list for accepting results from the invocation. - iree_vm_variant_list_t* outputs = nullptr; - IREE_ASSERT_OK( - iree_vm_variant_list_alloc(1, IREE_ALLOCATOR_SYSTEM, &outputs)); + vm::ref<iree_vm_list_t> outputs; + IREE_ASSERT_OK(iree_vm_list_create(/*element_type=*/nullptr, 1, + IREE_ALLOCATOR_SYSTEM, &outputs)); // Invoke the function. - IREE_ASSERT_OK(iree_vm_invoke( - context_, LookupFunction("string_tensor_to_string"), - /*policy=*/nullptr, inputs, outputs, IREE_ALLOCATOR_SYSTEM)); + IREE_ASSERT_OK(iree_vm_invoke(context_, + LookupFunction("string_tensor_to_string"), + /*policy=*/nullptr, inputs.get(), + outputs.get(), IREE_ALLOCATOR_SYSTEM)); - // Retrieve and validate the string tensor; - strings_string_t* output_string = - strings_string_deref(&iree_vm_variant_list_get(outputs, 0)->ref); + // Retrieve and validate the string tensor. + auto* output_string = + reinterpret_cast<strings_string_t*>(iree_vm_list_get_ref_deref( + outputs.get(), 0, strings_string_get_descriptor())); ASSERT_EQ(output_string->value.size, expected_output.length()); EXPECT_EQ( absl::string_view(output_string->value.data, output_string->value.size), expected_output); - - // Free the lists. - iree_vm_variant_list_free(inputs); - iree_vm_variant_list_free(outputs); } template <typename T, iree_hal_element_type_t E> @@ -179,33 +171,26 @@ CreateBufferView<T, E>(contents, shape, &input_buffer_view); // Construct the input list for execution. - iree_vm_variant_list_t* inputs = nullptr; - IREE_ASSERT_OK( - iree_vm_variant_list_alloc(1, IREE_ALLOCATOR_SYSTEM, &inputs)); + vm::ref<iree_vm_list_t> inputs; + IREE_ASSERT_OK(iree_vm_list_create(/*element_type=*/nullptr, 1, + IREE_ALLOCATOR_SYSTEM, &inputs)); // Add the buffer view to the input list. - iree_vm_ref_t input_buffer_view_ref = - iree_hal_buffer_view_move_ref(input_buffer_view.get()); - IREE_ASSERT_OK( - iree_vm_variant_list_append_ref_retain(inputs, &input_buffer_view_ref)); + iree_vm_list_push_ref_retain(inputs.get(), input_buffer_view)); // Construct the output list for accepting results from the invocation. - iree_vm_variant_list_t* outputs = nullptr; - IREE_ASSERT_OK( - iree_vm_variant_list_alloc(1, IREE_ALLOCATOR_SYSTEM, &outputs)); + vm::ref<iree_vm_list_t> outputs; + IREE_ASSERT_OK(iree_vm_list_create(/*element_type=*/nullptr, 1, + IREE_ALLOCATOR_SYSTEM, &outputs)); // Invoke the function. IREE_ASSERT_OK(iree_vm_invoke(context_, LookupFunction("to_string_tensor"), - /*policy=*/nullptr, inputs, outputs, - IREE_ALLOCATOR_SYSTEM)); + /*policy=*/nullptr, inputs.get(), + outputs.get(), IREE_ALLOCATOR_SYSTEM)); // Compare the output to the expected result. CompareResults(expected, shape, outputs); - - // Free the lists. - iree_vm_variant_list_free(inputs); - iree_vm_variant_list_free(outputs); } void TestGather(absl::Span<const iree_string_view_t> dict, @@ -219,43 +204,36 @@ dict_shape.size(), &dict_string_tensor)); // Construct the input list for execution. - iree_vm_variant_list_t* inputs = nullptr; - IREE_ASSERT_OK( - iree_vm_variant_list_alloc(2, IREE_ALLOCATOR_SYSTEM, &inputs)); + vm::ref<iree_vm_list_t> inputs; + IREE_ASSERT_OK(iree_vm_list_create(/*element_type=*/nullptr, 2, + IREE_ALLOCATOR_SYSTEM, &inputs)); // Add the dict to the input list. iree_vm_ref_t dict_string_tensor_ref = strings_string_tensor_move_ref(dict_string_tensor.get()); - IREE_ASSERT_OK(iree_vm_variant_list_append_ref_retain( - inputs, &dict_string_tensor_ref)); + IREE_ASSERT_OK( + iree_vm_list_push_ref_retain(inputs.get(), &dict_string_tensor_ref)); vm::ref<iree_hal_buffer_view_t> input_buffer_view; CreateBufferView<int32_t, IREE_HAL_ELEMENT_TYPE_SINT_32>( ids, ids_shape, &input_buffer_view); // Add the ids tensor to the input list. - iree_vm_ref_t input_buffer_view_ref = - iree_hal_buffer_view_move_ref(input_buffer_view.get()); - IREE_ASSERT_OK( - iree_vm_variant_list_append_ref_retain(inputs, &input_buffer_view_ref)); + iree_vm_list_push_ref_retain(inputs.get(), input_buffer_view)); // Construct the output list for accepting results from the invocation. - iree_vm_variant_list_t* outputs = nullptr; - IREE_ASSERT_OK( - iree_vm_variant_list_alloc(1, IREE_ALLOCATOR_SYSTEM, &outputs)); + vm::ref<iree_vm_list_t> outputs; + IREE_ASSERT_OK(iree_vm_list_create(/*element_type=*/nullptr, 1, + IREE_ALLOCATOR_SYSTEM, &outputs)); // Invoke the function. IREE_ASSERT_OK(iree_vm_invoke(context_, LookupFunction("gather"), - /*policy=*/nullptr, inputs, outputs, - IREE_ALLOCATOR_SYSTEM)); + /*policy=*/nullptr, inputs.get(), + outputs.get(), IREE_ALLOCATOR_SYSTEM)); // Compare the output to the expected result. - CompareResults(expected, ids_shape, outputs); - - // Free the lists. - iree_vm_variant_list_free(inputs); - iree_vm_variant_list_free(outputs); + CompareResults(expected, ids_shape, std::move(outputs)); } void TestConcat(absl::Span<const iree_string_view_t> string_views, @@ -267,43 +245,37 @@ shape.data(), shape.size(), &string_tensor)); // Construct the input list for execution. - iree_vm_variant_list_t* inputs = nullptr; - IREE_ASSERT_OK( - iree_vm_variant_list_alloc(1, IREE_ALLOCATOR_SYSTEM, &inputs)); + vm::ref<iree_vm_list_t> inputs; + IREE_ASSERT_OK(iree_vm_list_create(/*element_type=*/nullptr, 1, + IREE_ALLOCATOR_SYSTEM, &inputs)); // Add the dict to the input list. - iree_vm_ref_t string_tensor_ref = - strings_string_tensor_move_ref(string_tensor.get()); - IREE_ASSERT_OK( - iree_vm_variant_list_append_ref_retain(inputs, &string_tensor_ref)); + IREE_ASSERT_OK(iree_vm_list_push_ref_retain(inputs.get(), string_tensor)); // Construct the output list for accepting results from the invocation. - iree_vm_variant_list_t* outputs = nullptr; - IREE_ASSERT_OK( - iree_vm_variant_list_alloc(1, IREE_ALLOCATOR_SYSTEM, &outputs)); + vm::ref<iree_vm_list_t> outputs; + IREE_ASSERT_OK(iree_vm_list_create(/*element_type=*/nullptr, 1, + IREE_ALLOCATOR_SYSTEM, &outputs)); // Invoke the function. IREE_ASSERT_OK(iree_vm_invoke(context_, LookupFunction("concat"), - /*policy=*/nullptr, inputs, outputs, - IREE_ALLOCATOR_SYSTEM)); + /*policy=*/nullptr, inputs.get(), + outputs.get(), IREE_ALLOCATOR_SYSTEM)); // Remove the last dimension from the shape to get the expected shape shape.remove_suffix(1); // Compare the output to the expected result. - CompareResults(expected, shape, outputs); - - // Free the lists. - iree_vm_variant_list_free(inputs); - iree_vm_variant_list_free(outputs); + CompareResults(expected, shape, std::move(outputs)); } void CompareResults(absl::Span<const iree_string_view_t> expected, absl::Span<const int32_t> expected_shape, - iree_vm_variant_list_t* outputs) { - // Retrieve and validate the string tensor; - strings_string_tensor_t* output_tensor = - strings_string_tensor_deref(&iree_vm_variant_list_get(outputs, 0)->ref); + vm::ref<iree_vm_list_t> outputs) { + // Retrieve and validate the string tensor. + auto* output_tensor = + reinterpret_cast<strings_string_tensor_t*>(iree_vm_list_get_ref_deref( + outputs.get(), 0, strings_string_tensor_get_descriptor())); // Validate the count. size_t count; @@ -348,26 +320,24 @@ std::string expected_output = "42\n"; // Construct the input list for execution. - iree_vm_variant_list_t* inputs = nullptr; - IREE_ASSERT_OK(iree_vm_variant_list_alloc(1, IREE_ALLOCATOR_SYSTEM, &inputs)); + vm::ref<iree_vm_list_t> inputs; + IREE_ASSERT_OK(iree_vm_list_create(/*element_type=*/nullptr, 1, + IREE_ALLOCATOR_SYSTEM, &inputs)); // Add the value parameter. iree_vm_value_t value = iree_vm_value_make_i32(input); - IREE_ASSERT_OK(iree_vm_variant_list_append_value(inputs, value)); + IREE_ASSERT_OK(iree_vm_list_push_value(inputs.get(), &value)); // Prepare outputs list to accept the results from the invocation. - iree_vm_variant_list_t* outputs = nullptr; - IREE_ASSERT_OK( - iree_vm_variant_list_alloc(0, IREE_ALLOCATOR_SYSTEM, &outputs)); + vm::ref<iree_vm_list_t> outputs; + IREE_ASSERT_OK(iree_vm_list_create(/*element_type=*/nullptr, 0, + IREE_ALLOCATOR_SYSTEM, &outputs)); CaptureStdout(); IREE_ASSERT_OK(iree_vm_invoke(context_, LookupFunction("print_example_func"), - /*policy=*/nullptr, inputs, outputs, + /*policy=*/nullptr, inputs.get(), outputs.get(), IREE_ALLOCATOR_SYSTEM)); EXPECT_EQ(GetCapturedStdout(), expected_output); - - iree_vm_variant_list_free(inputs); - iree_vm_variant_list_free(outputs); } TEST_F(StringsModuleTest, StringTensorToString_Scalar) {
diff --git a/iree/modules/tensorlist/BUILD b/iree/modules/tensorlist/BUILD index 1bc2d81..4e7a784 100644 --- a/iree/modules/tensorlist/BUILD +++ b/iree/modules/tensorlist/BUILD
@@ -41,9 +41,7 @@ "//iree/testing:gtest_main", "//iree/vm", "//iree/vm:bytecode_module", - "//iree/vm:ref", "//iree/vm:ref_cc", - "//iree/vm:variant_list", "@com_google_absl//absl/base:core_headers", "@com_google_absl//absl/container:inlined_vector", "@com_google_absl//absl/strings",
diff --git a/iree/modules/tensorlist/CMakeLists.txt b/iree/modules/tensorlist/CMakeLists.txt index eb2880a..affd749 100644 --- a/iree/modules/tensorlist/CMakeLists.txt +++ b/iree/modules/tensorlist/CMakeLists.txt
@@ -46,9 +46,7 @@ iree::testing::gtest_main iree::vm iree::vm::bytecode_module - iree::vm::ref iree::vm::ref_cc - iree::vm::variant_list ) iree_cc_library(
diff --git a/iree/modules/tensorlist/native_module.cc b/iree/modules/tensorlist/native_module.cc index c3fb8dd..9ca5c0a 100644 --- a/iree/modules/tensorlist/native_module.cc +++ b/iree/modules/tensorlist/native_module.cc
@@ -322,12 +322,14 @@ static iree_vm_ref_type_descriptor_t iree_tensorlist_descriptor = {0}; // Register our type with the vm::ref<T> static machinery. +namespace vm { template <> -struct ::iree::vm::ref_type_descriptor<TensorList> { +struct ref_type_descriptor<TensorList> { static const iree_vm_ref_type_descriptor_t* get() { return &iree_tensorlist_descriptor; } }; +} // namespace vm extern "C" iree_status_t iree_tensorlist_module_register_types() { static bool has_registered = false;
diff --git a/iree/modules/tensorlist/tensorlist_test.cc b/iree/modules/tensorlist/tensorlist_test.cc index e254a35..76227ee 100644 --- a/iree/modules/tensorlist/tensorlist_test.cc +++ b/iree/modules/tensorlist/tensorlist_test.cc
@@ -27,9 +27,7 @@ #include "iree/testing/gtest.h" #include "iree/vm/api.h" #include "iree/vm/bytecode_module.h" -#include "iree/vm/ref.h" #include "iree/vm/ref_cc.h" -#include "iree/vm/variant_list.h" namespace iree { @@ -135,26 +133,27 @@ CreateBufferView(kBufferContents, shape, device_, &input_buffer_view); // Pass in the tensor as a HAL buffer view. - iree_vm_variant_list_t* inputs = nullptr; - IREE_ASSERT_OK(iree_vm_variant_list_alloc(1, IREE_ALLOCATOR_SYSTEM, &inputs)); + vm::ref<iree_vm_list_t> inputs; + IREE_ASSERT_OK(iree_vm_list_create(/*element_type=*/nullptr, 1, + IREE_ALLOCATOR_SYSTEM, &inputs)); iree_vm_ref_t input_buffer_view_ref = iree_hal_buffer_view_move_ref(input_buffer_view.get()); IREE_ASSERT_OK( - iree_vm_variant_list_append_ref_retain(inputs, &input_buffer_view_ref)); + iree_vm_list_push_ref_retain(inputs.get(), &input_buffer_view_ref)); // Prepare outputs list to accept the results from the invocation. - iree_vm_variant_list_t* outputs = nullptr; - IREE_ASSERT_OK( - iree_vm_variant_list_alloc(1, IREE_ALLOCATOR_SYSTEM, &outputs)); + vm::ref<iree_vm_list_t> outputs; + IREE_ASSERT_OK(iree_vm_list_create(/*element_type=*/nullptr, 1, + IREE_ALLOCATOR_SYSTEM, &outputs)); // Synchronously invoke the function. IREE_ASSERT_OK(iree_vm_invoke( context_, LookupFunction("identity_through_set_item_get_item"), - /*policy=*/nullptr, inputs, outputs, IREE_ALLOCATOR_SYSTEM)); - iree_vm_variant_list_free(inputs); + /*policy=*/nullptr, inputs.get(), outputs.get(), IREE_ALLOCATOR_SYSTEM)); - iree_hal_buffer_view_t* returned_buffer_view = - iree_hal_buffer_view_deref(&iree_vm_variant_list_get(outputs, 0)->ref); + auto* returned_buffer_view = + reinterpret_cast<iree_hal_buffer_view_t*>(iree_vm_list_get_ref_deref( + outputs.get(), 0, iree_hal_buffer_view_get_descriptor())); ASSERT_NE(nullptr, returned_buffer_view); iree_hal_buffer_t* returned_buffer = iree_hal_buffer_view_buffer(returned_buffer_view); @@ -167,8 +166,6 @@ EXPECT_EQ(reinterpret_cast<float*>(mapped_memory.contents.data)[0], kBufferContents[0]); IREE_ASSERT_OK(iree_hal_buffer_unmap(returned_buffer, &mapped_memory)); - - iree_vm_variant_list_free(outputs); } TEST_F(TensorListModulesTest, IdentityThroughConcat) { @@ -179,26 +176,27 @@ CreateBufferView(kBufferContents, shape, device_, &input_buffer_view); // Pass in the tensor as a HAL buffer view. - iree_vm_variant_list_t* inputs = nullptr; - IREE_ASSERT_OK(iree_vm_variant_list_alloc(1, IREE_ALLOCATOR_SYSTEM, &inputs)); + vm::ref<iree_vm_list_t> inputs; + IREE_ASSERT_OK(iree_vm_list_create(/*element_type=*/nullptr, 1, + IREE_ALLOCATOR_SYSTEM, &inputs)); iree_vm_ref_t input_buffer_view_ref = iree_hal_buffer_view_move_ref(input_buffer_view.get()); IREE_ASSERT_OK( - iree_vm_variant_list_append_ref_retain(inputs, &input_buffer_view_ref)); + iree_vm_list_push_ref_retain(inputs.get(), &input_buffer_view_ref)); // Prepare outputs list to accept the results from the invocation. - iree_vm_variant_list_t* outputs = nullptr; - IREE_ASSERT_OK( - iree_vm_variant_list_alloc(1, IREE_ALLOCATOR_SYSTEM, &outputs)); + vm::ref<iree_vm_list_t> outputs; + IREE_ASSERT_OK(iree_vm_list_create(/*element_type=*/nullptr, 1, + IREE_ALLOCATOR_SYSTEM, &outputs)); // Synchronously invoke the function. IREE_ASSERT_OK(iree_vm_invoke( context_, LookupFunction("identity_through_concat"), - /*policy=*/nullptr, inputs, outputs, IREE_ALLOCATOR_SYSTEM)); - iree_vm_variant_list_free(inputs); + /*policy=*/nullptr, inputs.get(), outputs.get(), IREE_ALLOCATOR_SYSTEM)); - iree_hal_buffer_view_t* returned_buffer_view = - iree_hal_buffer_view_deref(&iree_vm_variant_list_get(outputs, 0)->ref); + auto* returned_buffer_view = + reinterpret_cast<iree_hal_buffer_view_t*>(iree_vm_list_get_ref_deref( + outputs.get(), 0, iree_hal_buffer_view_get_descriptor())); ASSERT_NE(nullptr, returned_buffer_view); iree_hal_buffer_t* returned_buffer = iree_hal_buffer_view_buffer(returned_buffer_view); @@ -222,8 +220,6 @@ mapped_memory.contents.data_length), 0); IREE_ASSERT_OK(iree_hal_buffer_unmap(returned_buffer, &mapped_memory)); - - iree_vm_variant_list_free(outputs); } TEST_F(TensorListModulesTest, IdentityThroughStack) { @@ -234,26 +230,27 @@ CreateBufferView(kBufferContents, shape, device_, &input_buffer_view); // Pass in the tensor as a HAL buffer view. - iree_vm_variant_list_t* inputs = nullptr; - IREE_ASSERT_OK(iree_vm_variant_list_alloc(1, IREE_ALLOCATOR_SYSTEM, &inputs)); + vm::ref<iree_vm_list_t> inputs; + IREE_ASSERT_OK(iree_vm_list_create(/*element_type=*/nullptr, 1, + IREE_ALLOCATOR_SYSTEM, &inputs)); iree_vm_ref_t input_buffer_view_ref = iree_hal_buffer_view_move_ref(input_buffer_view.get()); IREE_ASSERT_OK( - iree_vm_variant_list_append_ref_retain(inputs, &input_buffer_view_ref)); + iree_vm_list_push_ref_retain(inputs.get(), &input_buffer_view_ref)); // Prepare outputs list to accept the results from the invocation. - iree_vm_variant_list_t* outputs = nullptr; - IREE_ASSERT_OK( - iree_vm_variant_list_alloc(1, IREE_ALLOCATOR_SYSTEM, &outputs)); + vm::ref<iree_vm_list_t> outputs; + IREE_ASSERT_OK(iree_vm_list_create(/*element_type=*/nullptr, 1, + IREE_ALLOCATOR_SYSTEM, &outputs)); // Synchronously invoke the function. IREE_ASSERT_OK(iree_vm_invoke( context_, LookupFunction("identity_through_stack"), - /*policy=*/nullptr, inputs, outputs, IREE_ALLOCATOR_SYSTEM)); - iree_vm_variant_list_free(inputs); + /*policy=*/nullptr, inputs.get(), outputs.get(), IREE_ALLOCATOR_SYSTEM)); - iree_hal_buffer_view_t* returned_buffer_view = - iree_hal_buffer_view_deref(&iree_vm_variant_list_get(outputs, 0)->ref); + auto* returned_buffer_view = + reinterpret_cast<iree_hal_buffer_view_t*>(iree_vm_list_get_ref_deref( + outputs.get(), 0, iree_hal_buffer_view_get_descriptor())); ASSERT_NE(nullptr, returned_buffer_view); iree_hal_buffer_t* returned_buffer = iree_hal_buffer_view_buffer(returned_buffer_view); @@ -268,8 +265,6 @@ mapped_memory.contents.data_length), 0); IREE_ASSERT_OK(iree_hal_buffer_unmap(returned_buffer, &mapped_memory)); - - iree_vm_variant_list_free(outputs); } } // namespace
diff --git a/iree/samples/custom_modules/BUILD b/iree/samples/custom_modules/BUILD index 7a65e51..036f5b2 100644 --- a/iree/samples/custom_modules/BUILD +++ b/iree/samples/custom_modules/BUILD
@@ -42,7 +42,7 @@ "//iree/testing:gtest_main", "//iree/vm", "//iree/vm:bytecode_module", - "//iree/vm:ref", + "//iree/vm:ref_cc", "@com_google_absl//absl/base:core_headers", "@com_google_absl//absl/strings", ],
diff --git a/iree/samples/custom_modules/CMakeLists.txt b/iree/samples/custom_modules/CMakeLists.txt index 44224a4..b98d0f4 100644 --- a/iree/samples/custom_modules/CMakeLists.txt +++ b/iree/samples/custom_modules/CMakeLists.txt
@@ -46,7 +46,7 @@ iree::testing::gtest_main iree::vm iree::vm::bytecode_module - iree::vm::ref + iree::vm::ref_cc ) iree_cc_library(
diff --git a/iree/samples/custom_modules/custom_modules_test.cc b/iree/samples/custom_modules/custom_modules_test.cc index 2bb38ca..51a0176 100644 --- a/iree/samples/custom_modules/custom_modules_test.cc +++ b/iree/samples/custom_modules/custom_modules_test.cc
@@ -25,7 +25,7 @@ #include "iree/testing/gtest.h" #include "iree/vm/api.h" #include "iree/vm/bytecode_module.h" -#include "iree/vm/ref.h" +#include "iree/vm/ref_cc.h" namespace { @@ -104,34 +104,32 @@ // Pass in the message and number of times to print it. // TODO(benvanik): make a macro/magic. - iree_vm_variant_list_t* inputs = nullptr; - IREE_ASSERT_OK(iree_vm_variant_list_alloc(2, IREE_ALLOCATOR_SYSTEM, &inputs)); + iree::vm::ref<iree_vm_list_t> inputs; + IREE_ASSERT_OK(iree_vm_list_create(/*element_type=*/nullptr, 2, + IREE_ALLOCATOR_SYSTEM, &inputs)); iree_vm_ref_t input_message_ref = iree_custom_message_move_ref(input_message); - IREE_ASSERT_OK( - iree_vm_variant_list_append_ref_move(inputs, &input_message_ref)); - IREE_ASSERT_OK(iree_vm_variant_list_append_value(inputs, count)); + IREE_ASSERT_OK(iree_vm_list_push_ref_move(inputs.get(), &input_message_ref)); + IREE_ASSERT_OK(iree_vm_list_push_value(inputs.get(), &count)); // Prepare outputs list to accept the results from the invocation. - iree_vm_variant_list_t* outputs = nullptr; - IREE_ASSERT_OK( - iree_vm_variant_list_alloc(1, IREE_ALLOCATOR_SYSTEM, &outputs)); + iree::vm::ref<iree_vm_list_t> outputs; + IREE_ASSERT_OK(iree_vm_list_create(/*element_type=*/nullptr, 1, + IREE_ALLOCATOR_SYSTEM, &outputs)); // Synchronously invoke the function. IREE_ASSERT_OK(iree_vm_invoke(context_, LookupFunction("reverseAndPrint"), - /*policy=*/nullptr, inputs, outputs, + /*policy=*/nullptr, inputs.get(), outputs.get(), IREE_ALLOCATOR_SYSTEM)); - iree_vm_variant_list_free(inputs); // Read back the message that we reversed inside of the module. iree_custom_message_t* reversed_message = - iree_custom_message_deref(&iree_vm_variant_list_get(outputs, 0)->ref); + (iree_custom_message_t*)iree_vm_list_get_ref_deref( + outputs.get(), 0, iree_custom_message_get_descriptor()); ASSERT_NE(nullptr, reversed_message); char result_buffer[256]; IREE_ASSERT_OK(iree_custom_message_read_value(reversed_message, result_buffer, ABSL_ARRAYSIZE(result_buffer))); EXPECT_STREQ("!dlrow olleh", result_buffer); - - iree_vm_variant_list_free(outputs); } TEST_F(CustomModulesTest, PrintTensor) { @@ -147,33 +145,31 @@ IREE_ALLOCATOR_SYSTEM, IREE_ALLOCATOR_SYSTEM, &buffer)); // Pass in the tensor as an expanded HAL buffer. - iree_vm_variant_list_t* inputs = nullptr; - IREE_ASSERT_OK(iree_vm_variant_list_alloc(1, IREE_ALLOCATOR_SYSTEM, &inputs)); + iree::vm::ref<iree_vm_list_t> inputs; + IREE_ASSERT_OK(iree_vm_list_create(/*element_type=*/nullptr, 1, + IREE_ALLOCATOR_SYSTEM, &inputs)); iree_vm_ref_t input_buffer_ref = iree_hal_buffer_move_ref(buffer); - IREE_ASSERT_OK( - iree_vm_variant_list_append_ref_move(inputs, &input_buffer_ref)); + IREE_ASSERT_OK(iree_vm_list_push_ref_move(inputs.get(), &input_buffer_ref)); // Prepare outputs list to accept the results from the invocation. - iree_vm_variant_list_t* outputs = nullptr; - IREE_ASSERT_OK( - iree_vm_variant_list_alloc(1, IREE_ALLOCATOR_SYSTEM, &outputs)); + iree::vm::ref<iree_vm_list_t> outputs; + IREE_ASSERT_OK(iree_vm_list_create(/*element_type=*/nullptr, 1, + IREE_ALLOCATOR_SYSTEM, &outputs)); // Synchronously invoke the function. IREE_ASSERT_OK(iree_vm_invoke(context_, LookupFunction("printTensor"), - /*policy=*/nullptr, inputs, outputs, + /*policy=*/nullptr, inputs.get(), outputs.get(), IREE_ALLOCATOR_SYSTEM)); - iree_vm_variant_list_free(inputs); // Read back the message that we printed inside of the module. iree_custom_message_t* printed_message = - iree_custom_message_deref(&iree_vm_variant_list_get(outputs, 0)->ref); + (iree_custom_message_t*)iree_vm_list_get_ref_deref( + outputs.get(), 0, iree_custom_message_get_descriptor()); ASSERT_NE(nullptr, printed_message); char result_buffer[256]; IREE_ASSERT_OK(iree_custom_message_read_value(printed_message, result_buffer, ABSL_ARRAYSIZE(result_buffer))); EXPECT_STREQ("2x4xf32=[0 1 2 3][4 5 6 7]", result_buffer); - - iree_vm_variant_list_free(outputs); } TEST_F(CustomModulesTest, RoundTripTensor) { @@ -189,33 +185,31 @@ IREE_ALLOCATOR_SYSTEM, IREE_ALLOCATOR_SYSTEM, &buffer)); // Pass in the tensor as an expanded HAL buffer. - iree_vm_variant_list_t* inputs = nullptr; - IREE_ASSERT_OK(iree_vm_variant_list_alloc(1, IREE_ALLOCATOR_SYSTEM, &inputs)); + iree::vm::ref<iree_vm_list_t> inputs; + IREE_ASSERT_OK(iree_vm_list_create(/*element_type=*/nullptr, 1, + IREE_ALLOCATOR_SYSTEM, &inputs)); iree_vm_ref_t input_buffer_ref = iree_hal_buffer_move_ref(buffer); - IREE_ASSERT_OK( - iree_vm_variant_list_append_ref_move(inputs, &input_buffer_ref)); + IREE_ASSERT_OK(iree_vm_list_push_ref_move(inputs.get(), &input_buffer_ref)); // Prepare outputs list to accept the results from the invocation. - iree_vm_variant_list_t* outputs = nullptr; - IREE_ASSERT_OK( - iree_vm_variant_list_alloc(1, IREE_ALLOCATOR_SYSTEM, &outputs)); + iree::vm::ref<iree_vm_list_t> outputs; + IREE_ASSERT_OK(iree_vm_list_create(/*element_type=*/nullptr, 1, + IREE_ALLOCATOR_SYSTEM, &outputs)); // Synchronously invoke the function. IREE_ASSERT_OK(iree_vm_invoke(context_, LookupFunction("roundTripTensor"), - /*policy=*/nullptr, inputs, outputs, + /*policy=*/nullptr, inputs.get(), outputs.get(), IREE_ALLOCATOR_SYSTEM)); - iree_vm_variant_list_free(inputs); // Read back the message that's been moved around. iree_custom_message_t* printed_message = - iree_custom_message_deref(&iree_vm_variant_list_get(outputs, 0)->ref); + (iree_custom_message_t*)iree_vm_list_get_ref_deref( + outputs.get(), 0, iree_custom_message_get_descriptor()); ASSERT_NE(nullptr, printed_message); char result_buffer[256]; IREE_ASSERT_OK(iree_custom_message_read_value(printed_message, result_buffer, ABSL_ARRAYSIZE(result_buffer))); EXPECT_STREQ("2x4xf32=[0 1 2 3][4 5 6 7]", result_buffer); - - iree_vm_variant_list_free(outputs); } } // namespace
diff --git a/iree/samples/simple_embedding/BUILD b/iree/samples/simple_embedding/BUILD index 7a9ad60..5dad5bf 100644 --- a/iree/samples/simple_embedding/BUILD +++ b/iree/samples/simple_embedding/BUILD
@@ -49,8 +49,9 @@ "//iree/base:logging", "//iree/hal:api", "//iree/modules/hal", - "//iree/vm:bytecode_module", "//iree/vm", + "//iree/vm:bytecode_module", + "//iree/vm:ref_cc", # These are the drivers we support running with and can produce # executables for from the source MLIR.
diff --git a/iree/samples/simple_embedding/CMakeLists.txt b/iree/samples/simple_embedding/CMakeLists.txt index 2e648e5..1001ec8 100644 --- a/iree/samples/simple_embedding/CMakeLists.txt +++ b/iree/samples/simple_embedding/CMakeLists.txt
@@ -49,6 +49,7 @@ iree::testing::gtest_main iree::vm iree::vm::bytecode_module + iree::vm::ref_cc LABELS "driver=vulkan" )
diff --git a/iree/samples/simple_embedding/simple_embedding_test.cc b/iree/samples/simple_embedding/simple_embedding_test.cc index 4e4cc9a..489ac33 100644 --- a/iree/samples/simple_embedding/simple_embedding_test.cc +++ b/iree/samples/simple_embedding/simple_embedding_test.cc
@@ -23,6 +23,7 @@ #include "iree/testing/gtest.h" #include "iree/vm/api.h" #include "iree/vm/bytecode_module.h" +#include "iree/vm/ref_cc.h" // Compiled module embedded here to avoid file IO: #include "iree/samples/simple_embedding/simple_embedding_test_bytecode_module.h" @@ -144,31 +145,30 @@ // Setup call inputs with our buffers. // TODO(benvanik): make a macro/magic. - iree_vm_variant_list_t* inputs = nullptr; - IREE_ASSERT_OK(iree_vm_variant_list_alloc(2, IREE_ALLOCATOR_SYSTEM, &inputs)); + vm::ref<iree_vm_list_t> inputs; + IREE_ASSERT_OK(iree_vm_list_create(/*element_type=*/nullptr, 2, + IREE_ALLOCATOR_SYSTEM, &inputs)); auto arg0_buffer_ref = iree_hal_buffer_move_ref(arg0_buffer); auto arg1_buffer_ref = iree_hal_buffer_move_ref(arg1_buffer); - IREE_ASSERT_OK( - iree_vm_variant_list_append_ref_move(inputs, &arg0_buffer_ref)); - IREE_ASSERT_OK( - iree_vm_variant_list_append_ref_move(inputs, &arg1_buffer_ref)); + IREE_ASSERT_OK(iree_vm_list_push_ref_move(inputs.get(), &arg0_buffer_ref)); + IREE_ASSERT_OK(iree_vm_list_push_ref_move(inputs.get(), &arg1_buffer_ref)); // Prepare outputs list to accept the results from the invocation. - iree_vm_variant_list_t* outputs = nullptr; - IREE_ASSERT_OK( - iree_vm_variant_list_alloc(1, IREE_ALLOCATOR_SYSTEM, &outputs)); + vm::ref<iree_vm_list_t> outputs; + IREE_ASSERT_OK(iree_vm_list_create(/*element_type=*/nullptr, 1, + IREE_ALLOCATOR_SYSTEM, &outputs)); // Synchronously invoke the function. LOG(INFO) << "Calling " << kMainFunctionName << "..."; IREE_ASSERT_OK(iree_vm_invoke(context, main_function, - /*policy=*/nullptr, inputs, outputs, + /*policy=*/nullptr, inputs.get(), outputs.get(), IREE_ALLOCATOR_SYSTEM)); - iree_vm_variant_list_free(inputs); // Get the result buffers from the invocation. LOG(INFO) << "Retreiving results..."; - iree_hal_buffer_t* ret_buffer = - iree_hal_buffer_deref(&iree_vm_variant_list_get(outputs, 0)->ref); + auto* ret_buffer = + reinterpret_cast<iree_hal_buffer_t*>(iree_vm_list_get_ref_deref( + outputs.get(), 0, iree_hal_buffer_get_descriptor())); ASSERT_NE(nullptr, ret_buffer); // Read back the results and ensure we got the right values. @@ -183,8 +183,8 @@ ASSERT_API_OK(iree_hal_buffer_unmap(ret_buffer, &mapped_memory)); LOG(INFO) << "Results match!"; - iree_vm_variant_list_free(outputs); - + inputs.reset(); + outputs.reset(); iree_hal_device_release(device); iree_vm_context_release(context); iree_vm_instance_release(instance);
diff --git a/iree/samples/vulkan/BUILD b/iree/samples/vulkan/BUILD index 7bf6f7e..b684de7 100644 --- a/iree/samples/vulkan/BUILD +++ b/iree/samples/vulkan/BUILD
@@ -55,6 +55,7 @@ "//iree/modules/hal", "//iree/vm", "//iree/vm:bytecode_module", + "//iree/vm:ref_cc", "@com_google_absl//absl/base:core_headers", "@dear_imgui", "@dear_imgui//:imgui_sdl_vulkan",
diff --git a/iree/samples/vulkan/CMakeLists.txt b/iree/samples/vulkan/CMakeLists.txt index 17c7579..8d83188 100644 --- a/iree/samples/vulkan/CMakeLists.txt +++ b/iree/samples/vulkan/CMakeLists.txt
@@ -59,6 +59,7 @@ iree::samples::vulkan::simple_mul_bytecode_module_cc iree::vm iree::vm::bytecode_module + iree::vm::ref_cc SDL2-static Vulkan::Vulkan LINKOPTS
diff --git a/iree/samples/vulkan/vulkan_inference_gui.cc b/iree/samples/vulkan/vulkan_inference_gui.cc index b4b75ce..d939a76 100644 --- a/iree/samples/vulkan/vulkan_inference_gui.cc +++ b/iree/samples/vulkan/vulkan_inference_gui.cc
@@ -35,6 +35,7 @@ #include "iree/modules/hal/hal_module.h" #include "iree/vm/api.h" #include "iree/vm/bytecode_module.h" +#include "iree/vm/ref_cc.h" // Other dependencies (helpers, etc.) #include "absl/base/macros.h" @@ -723,36 +724,34 @@ iree_hal_buffer_release(input0_buffer); iree_hal_buffer_release(input1_buffer); // Marshal input buffer views through a VM variant list. - iree_vm_variant_list_t* input_list = nullptr; - CHECK_IREE_OK( - iree_vm_variant_list_alloc(2, IREE_ALLOCATOR_SYSTEM, &input_list)); + vm::ref<iree_vm_list_t> inputs; + CHECK_IREE_OK(iree_vm_list_create(/*element_type=*/nullptr, 2, + IREE_ALLOCATOR_SYSTEM, &inputs)); auto input0_buffer_view_ref = iree_hal_buffer_view_move_ref(input0_buffer_view); auto input1_buffer_view_ref = iree_hal_buffer_view_move_ref(input1_buffer_view); - CHECK_IREE_OK(iree_vm_variant_list_append_ref_move( - input_list, &input0_buffer_view_ref)); - CHECK_IREE_OK(iree_vm_variant_list_append_ref_move( - input_list, &input1_buffer_view_ref)); + CHECK_IREE_OK( + iree_vm_list_push_ref_move(inputs.get(), &input0_buffer_view_ref)); + CHECK_IREE_OK( + iree_vm_list_push_ref_move(inputs.get(), &input1_buffer_view_ref)); // Prepare outputs list to accept results from the invocation. - iree_vm_variant_list_t* output_list = nullptr; - CHECK_IREE_OK(iree_vm_variant_list_alloc(kElementCount * sizeof(float), - IREE_ALLOCATOR_SYSTEM, - &output_list)); + vm::ref<iree_vm_list_t> outputs; + CHECK_IREE_OK(iree_vm_list_create(/*element_type=*/nullptr, + kElementCount * sizeof(float), + IREE_ALLOCATOR_SYSTEM, &outputs)); // Synchronously invoke the function. CHECK_IREE_OK(iree_vm_invoke(iree_context, main_function, - /*policy=*/nullptr, input_list, - output_list, IREE_ALLOCATOR_SYSTEM)); - iree_vm_variant_list_free(input_list); + /*policy=*/nullptr, inputs.get(), + outputs.get(), IREE_ALLOCATOR_SYSTEM)); // Read back the results. DLOG(INFO) << "Reading back results..."; - iree_vm_variant_t* output_variant = - iree_vm_variant_list_get(output_list, 0); - auto* output_buffer_view = - iree_hal_buffer_view_deref(&output_variant->ref); + auto* output_buffer_view = reinterpret_cast<iree_hal_buffer_view_t*>( + iree_vm_list_get_ref_deref(outputs.get(), 0, + iree_hal_buffer_view_get_descriptor())); auto* output_buffer = iree_hal_buffer_view_buffer(output_buffer_view); iree_hal_mapped_memory_t mapped_memory; CHECK_IREE_OK(iree_hal_buffer_map(output_buffer, @@ -761,7 +760,6 @@ memcpy(&latest_output, mapped_memory.contents.data, mapped_memory.contents.data_length); iree_hal_buffer_unmap(output_buffer, &mapped_memory); - iree_vm_variant_list_free(output_list); dirty = false; }
diff --git a/iree/schemas/bytecode_module_def.fbs b/iree/schemas/bytecode_module_def.fbs index f431332..2394afb 100644 --- a/iree/schemas/bytecode_module_def.fbs +++ b/iree/schemas/bytecode_module_def.fbs
@@ -43,7 +43,7 @@ // Function level reflection attributes. // These are typically used to communicate additional ABI metadata needed // for dynamic invocation and host language mapping. - // See: docs/function_abi.md + // See: docs/design_docs/function_abi.md reflection_attrs:[ReflectionAttrDef]; }
diff --git a/iree/test/e2e/regression/dynamic_torch_index_select_high_rank.mlir b/iree/test/e2e/regression/dynamic_torch_index_select_high_rank.mlir new file mode 100644 index 0000000..ce42e31 --- /dev/null +++ b/iree/test/e2e/regression/dynamic_torch_index_select_high_rank.mlir
@@ -0,0 +1,58 @@ +// RUN: iree-run-mlir %s -iree-hal-target-backends=llvm-ir -input-value="2x2xi32=[6, 7] [8, 9]" -input-value="2x2x2x2xi32=[[[0, 1] [1, 0]] [[0, 0] [1, 1]]] [[[1, 1] [0, 0]] [[0, 1] [1, 0]]]" | IreeFileCheck %s + +// CHECK-LABEL: EXEC @torch_index_select1 +func @torch_index_select1(%arg0: tensor<?x?xi32>, %arg1: tensor<?x?x?x?xi32>) -> tensor<?x?x?x?xi32> attributes {iree.module.export} { + %0 = "mhlo.torch_index_select"(%arg0, %arg1) {batch_dims = 1 : i64, dim = 1 : i64} : (tensor<?x?xi32>, tensor<?x?x?x?xi32>) -> tensor<?x?x?x?xi32> + return %0 : tensor<?x?x?x?xi32> +} + +// CHECK: 2x2x2x2xi32=[ +// CHECK-SAME: [ +// CHECK-SAME: [6 7][7 6] +// CHECK-SAME: ][ +// CHECK-SAME: [6 6][7 7] +// CHECK-SAME: ] +// CHECK-SAME: ][ +// CHECK-SAME: [ +// CHECK-SAME: [9 9][8 8] +// CHECK-SAME: ][ +// CHECK-SAME: [8 9][9 8] +// CHECK-SAME: ] +// CHECK-SAME: ] + +// CHECK-LABEL: EXEC @torch_index_select2 +func @torch_index_select2(%arg0: tensor<?x?xi32>, %arg1: tensor<?x?x?x?xi32>) -> tensor<?x?x?x?x?xi32> attributes {iree.module.export} { + %0 = "mhlo.torch_index_select"(%arg0, %arg1) {batch_dims = 0 : i64, dim = 0 : i64} : (tensor<?x?xi32>, tensor<?x?x?x?xi32>) -> tensor<?x?x?x?x?xi32> + return %0 : tensor<?x?x?x?x?xi32> +} + +// CHECK: 2x2x2x2x2xi32=[ +// CHECK-SAME: [ +// CHECK-SAME: [ +// CHECK-SAME: [6 7][8 9] +// CHECK-SAME: ][ +// CHECK-SAME: [8 9][6 7] +// CHECK-SAME: ] +// CHECK-SAME: ][ +// CHECK-SAME: [ +// CHECK-SAME: [6 7][6 7] +// CHECK-SAME: ][ +// CHECK-SAME: [8 9][8 9] +// CHECK-SAME: ] +// CHECK-SAME: ] +// CHECK-SAME: ][ +// CHECK-SAME: [ +// CHECK-SAME: [ +// CHECK-SAME: [8 9][8 9] +// CHECK-SAME: ][ +// CHECK-SAME: [6 7][6 7] +// CHECK-SAME: ] +// CHECK-SAME: ][ +// CHECK-SAME: [ +// CHECK-SAME: [6 7][8 9] +// CHECK-SAME: ][ +// CHECK-SAME: [8 9][6 7] +// CHECK-SAME: ] +// CHECK-SAME: ] +// CHECK-SAME: ] +
diff --git a/iree/test/e2e/regression/dynamic_torch_index_select_negative.mlir b/iree/test/e2e/regression/dynamic_torch_index_select_negative.mlir new file mode 100644 index 0000000..0755b5c --- /dev/null +++ b/iree/test/e2e/regression/dynamic_torch_index_select_negative.mlir
@@ -0,0 +1,14 @@ +// RUN: iree-run-mlir %s -iree-hal-target-backends=llvm-ir -input-value="2x2x2xi32=[[100, 101] [110, 111]] [[200, 201] [210, 211]]" -input-value="2x2x2xi32=[[0, 1] [1, 0]] [[0, 0] [1, 1]]" | IreeFileCheck %s + +// CHECK-LABEL: EXEC @torch_index_select1 +func @torch_index_select1(%arg0: tensor<?x?x?xi32>, %arg1: tensor<?x?x?xi32>) -> tensor<?x?x?xi32> attributes {iree.module.export} { + %0 = "mhlo.torch_index_select"(%arg0, %arg1) {batch_dims = -1 : i64, dim = -1 : i64} : (tensor<?x?x?xi32>, tensor<?x?x?xi32>) -> tensor<?x?x?xi32> + return %0 : tensor<?x?x?xi32> +} + +// CHECK: 2x2x2xi32=[ +// CHECK-SAME: [100 101][111 110] +// CHECK-SAME: ][ +// CHECK-SAME: [200 200][211 211] +// CHECK-SAME: ] +
diff --git a/iree/test/e2e/regression/dynamic_torch_index_select_scalar.mlir b/iree/test/e2e/regression/dynamic_torch_index_select_scalar.mlir new file mode 100644 index 0000000..8ccb4fa --- /dev/null +++ b/iree/test/e2e/regression/dynamic_torch_index_select_scalar.mlir
@@ -0,0 +1,17 @@ +// RUN: iree-run-mlir %s -iree-hal-target-backends=llvm-ir -input-value="5x1x5xi32=[[1,2,3,4,5]] [[6,7,8,9,10]] [[11,12,13,14,15]] [[16,17,18,19,20]] [[21,22,23,24,25]]" -input-value="i32=0" | IreeFileCheck %s + +// CHECK-LABEL: EXEC @torch_index_select1 +func @torch_index_select1(%arg0: tensor<?x?x?xi32>, %arg1: tensor<i32>) -> tensor<?x?xi32> attributes {iree.module.export} { + %0 = "mhlo.torch_index_select"(%arg0, %arg1) {batch_dims = 0 : i64, dim = 0 : i64} : (tensor<?x?x?xi32>, tensor<i32>) -> tensor<?x?xi32> + return %0 : tensor<?x?xi32> +} + +// CHECK: 1x5xi32=[1 2 3 4 5] + +// CHECK-LABEL: EXEC @torch_index_select2 +func @torch_index_select2(%arg0: tensor<?x?x?xi32>, %arg1: tensor<i32>) -> tensor<?x?xi32> attributes {iree.module.export} { + %0 = "mhlo.torch_index_select"(%arg0, %arg1) {batch_dims = 0 : i64, dim = 1 : i64} : (tensor<?x?x?xi32>, tensor<i32>) -> tensor<?x?xi32> + return %0 : tensor<?x?xi32> +} + +// CHECK: 5x5xi32=[1 2 3 4 5][6 7 8 9 10][11 12 13 14 15][16 17 18 19 20][21 22 23 24 25]
diff --git a/iree/test/e2e/regression/dynamic_torch_index_select_vector.mlir b/iree/test/e2e/regression/dynamic_torch_index_select_vector.mlir new file mode 100644 index 0000000..d22ccc7 --- /dev/null +++ b/iree/test/e2e/regression/dynamic_torch_index_select_vector.mlir
@@ -0,0 +1,28 @@ +// RUN: iree-run-mlir %s -iree-hal-target-backends=llvm-ir -input-value="3x2x2xi32=[[1, 2] [3, 4]] [[5, 6] [7, 8]] [[9, 10] [11, 12]]" -input-value="2xi32=[0, 1]" | IreeFileCheck %s + +// CHECK-LABEL: EXEC @torch_index_select1 +func @torch_index_select1(%arg0: tensor<?x?x?xi32>, %arg1: tensor<?xi32>) -> tensor<?x?x?xi32> attributes {iree.module.export} { + %0 = "mhlo.torch_index_select"(%arg0, %arg1) {batch_dims = 0 : i64, dim = 1 : i64} : (tensor<?x?x?xi32>, tensor<?xi32>) -> tensor<?x?x?xi32> + return %0 : tensor<?x?x?xi32> +} + +// CHECK: 3x2x2xi32=[ +// CHECK-SAME: [1 2][3 4] +// CHECK-SAME: ][ +// CHECK-SAME: [5 6][7 8] +// CHECK-SAME: ][ +// CHECK-SAME: [9 10][11 12] +// CHECK-SAME: ] + +// CHECK-LABEL: EXEC @torch_index_select2 +func @torch_index_select2(%arg0: tensor<?x?x?xi32>, %arg1: tensor<?xi32>) -> tensor<?x?x?xi32> attributes {iree.module.export} { + %0 = "mhlo.torch_index_select"(%arg0, %arg1) {batch_dims = 0 : i64, dim = 0 : i64} : (tensor<?x?x?xi32>, tensor<?xi32>) -> tensor<?x?x?xi32> + return %0 : tensor<?x?x?xi32> +} + +// CHECK: 2x2x2xi32=[ +// CHECK-SAME: [1 2][3 4] +// CHECK-SAME: ][ +// CHECK-SAME: [5 6][7 8] +// CHECK-SAME: ] +
diff --git a/iree/test/e2e/vulkan_specific/BUILD b/iree/test/e2e/vulkan_specific/BUILD index 0565a86..4d20c33 100644 --- a/iree/test/e2e/vulkan_specific/BUILD +++ b/iree/test/e2e/vulkan_specific/BUILD
@@ -29,38 +29,3 @@ driver = "vulkan", target_backend = "vulkan-spirv", ) - -# TODO(#2345): Merge two tests into one single file. -iree_check_single_backend_test_suite( - name = "check_vulkan-spirv-split-pad-conv_vulkan", - srcs = [ - "convolution1.mlir", - "convolution2.mlir", - ], - driver = "vulkan", - target_backend = "vulkan-spirv", -) - -# TODO(#2345): Merge two tests into one single file. -iree_check_single_backend_test_suite( - name = "check_vulkan-spirv-nosplit-pad-conv_vulkan", - srcs = [ - "convolution1.mlir", - "convolution2.mlir", - ], - compiler_flags = ["-iree-extract-pad-from-conv=false"], - driver = "vulkan", - target_backend = "vulkan-spirv", -) - -# TODO(#2345): Merge two tests into one single file. -iree_check_single_backend_test_suite( - name = "check_vulkan-spirv-conv-nocontrol_vulkan", - srcs = [ - "convolution1.mlir", - "convolution2.mlir", - ], - compiler_flags = ["-iree-codegen-use-legacy-conv-lowering=false"], - driver = "vulkan", - target_backend = "vulkan-spirv", -)
diff --git a/iree/test/e2e/vulkan_specific/CMakeLists.txt b/iree/test/e2e/vulkan_specific/CMakeLists.txt index cca6c58..32ee021 100644 --- a/iree/test/e2e/vulkan_specific/CMakeLists.txt +++ b/iree/test/e2e/vulkan_specific/CMakeLists.txt
@@ -25,43 +25,3 @@ DRIVER vulkan ) - -iree_check_single_backend_test_suite( - NAME - check_vulkan-spirv-split-pad-conv_vulkan - SRCS - "convolution1.mlir" - "convolution2.mlir" - TARGET_BACKEND - vulkan-spirv - DRIVER - vulkan -) - -iree_check_single_backend_test_suite( - NAME - check_vulkan-spirv-nosplit-pad-conv_vulkan - SRCS - "convolution1.mlir" - "convolution2.mlir" - TARGET_BACKEND - vulkan-spirv - DRIVER - vulkan - COMPILER_FLAGS - "-iree-extract-pad-from-conv=false" -) - -iree_check_single_backend_test_suite( - NAME - check_vulkan-spirv-conv-nocontrol_vulkan - SRCS - "convolution1.mlir" - "convolution2.mlir" - TARGET_BACKEND - vulkan-spirv - DRIVER - vulkan - COMPILER_FLAGS - "-iree-codegen-use-legacy-conv-lowering=false" -)
diff --git a/iree/test/e2e/vulkan_specific/convolution1.mlir b/iree/test/e2e/vulkan_specific/convolution1.mlir deleted file mode 100644 index d0fc606..0000000 --- a/iree/test/e2e/vulkan_specific/convolution1.mlir +++ /dev/null
@@ -1,66 +0,0 @@ -func @conv2d_nopadding() attributes { iree.module.export } { - %inputs = iree.unfoldable_constant dense<[[ - [[ 1.0, 2.0], [ 3.0, 4.0], [ 5.0, 6.0], [ 7.0, 8.0], [ 9.0, 10.0]], - [[11.0, 12.0], [13.0, 14.0], [15.0, 16.0], [17.0, 18.0], [19.0, 20.0]], - [[21.0, 22.0], [23.0, 24.0], [25.0, 26.0], [27.0, 28.0], [29.0, 30.0]], - [[31.0, 32.0], [33.0, 34.0], [35.0, 36.0], [37.0, 38.0], [39.0, 40.0]]]]> : tensor<1x4x5x2xf32> - %weights = iree.unfoldable_constant dense<[ - [[[ 1.0], [ 2.0]], [[ 3.0], [ 4.0]]], - [[[ 5.0], [ 6.0]], [[ 7.0], [ 8.0]]], - [[[ 9.0], [10.0]], [[11.0], [12.0]]]]> : tensor<3x2x2x1xf32> - %res = "mhlo.convolution"(%inputs, %weights) { - batch_group_count = 1 : i64, - dimension_numbers = { - input_batch_dimension = 0 : i64, - input_feature_dimension = 3 : i64, - input_spatial_dimensions = dense<[1, 2]> : tensor<2xi64>, - kernel_input_feature_dimension = 2 : i64, - kernel_output_feature_dimension = 3 : i64, - kernel_spatial_dimensions = dense<[0, 1]> : tensor<2xi64>, - output_batch_dimension = 0 : i64, - output_feature_dimension = 3 : i64, - output_spatial_dimensions = dense<[1, 2]> : tensor<2xi64>}, - feature_group_count = 1 : i64, - rhs_dilation = dense<1> : tensor<2xi64>, - window_strides = dense<1> : tensor<2xi64>} : (tensor<1x4x5x2xf32>, tensor<3x2x2x1xf32>) -> tensor<1x2x3x1xf32> - check.expect_almost_eq_const(%res, dense<[[ - [[1310.0],[1466.0],[1622.0]], - [[2090.0],[2246.0],[2402.0]] - ]]> : tensor<1x2x3x1xf32>) : tensor<1x2x3x1xf32> - return -} - -func @conv2d_1452x3221_same() attributes { iree.module.export } { - %inputs = iree.unfoldable_constant dense<[[ - [[ 1.0, 2.0], [ 3.0, 4.0], [ 5.0, 6.0], [ 7.0, 8.0], [ 9.0, 10.0]], - [[11.0, 12.0], [13.0, 14.0], [15.0, 16.0], [17.0, 18.0], [19.0, 20.0]], - [[21.0, 22.0], [23.0, 24.0], [25.0, 26.0], [27.0, 28.0], [29.0, 30.0]], - [[31.0, 32.0], [33.0, 34.0], [35.0, 36.0], [37.0, 38.0], [39.0, 40.0]]]]> : tensor<1x4x5x2xf32> - %weights = iree.unfoldable_constant dense<[ - [[[ 1.0], [ 2.0]], [[ 3.0], [ 4.0]]], - [[[ 5.0], [ 6.0]], [[ 7.0], [ 8.0]]], - [[[ 9.0], [10.0]], [[11.0], [12.0]]]]> : tensor<3x2x2x1xf32> - %res = "mhlo.convolution"(%inputs, %weights) { - batch_group_count = 1 : i64, - dimension_numbers = { - input_batch_dimension = 0 : i64, - input_feature_dimension = 3 : i64, - input_spatial_dimensions = dense<[1, 2]> : tensor<2xi64>, - kernel_input_feature_dimension = 2 : i64, - kernel_output_feature_dimension = 3 : i64, - kernel_spatial_dimensions = dense<[0, 1]> : tensor<2xi64>, - output_batch_dimension = 0 : i64, - output_feature_dimension = 3 : i64, - output_spatial_dimensions = dense<[1, 2]> : tensor<2xi64>}, - feature_group_count = 1 : i64, - padding = dense<[[1, 1], [0, 1]]> : tensor<2x2xi64>, - rhs_dilation = dense<1> : tensor<2xi64>, - window_strides = dense<1> : tensor<2xi64>} : - (tensor<1x4x5x2xf32>, tensor<3x2x2x1xf32>) -> tensor<1x4x5x1xf32> - check.expect_almost_eq_const(%res, dense<[[ - [[ 600.0], [ 736.0], [ 872.0], [1008.0], [ 476.0]], - [[1310.0], [1466.0], [1622.0], [1778.0], [ 805.0]], - [[2090.0], [2246.0], [2402.0], [2558.0], [1135.0]], - [[1080.0], [1152.0], [1224.0], [1296.0], [ 524.0]]]]> : tensor<1x4x5x1xf32>) : tensor<1x4x5x1xf32> - return -}
diff --git a/iree/test/e2e/vulkan_specific/convolution2.mlir b/iree/test/e2e/vulkan_specific/convolution2.mlir deleted file mode 100644 index ce88d5d..0000000 --- a/iree/test/e2e/vulkan_specific/convolution2.mlir +++ /dev/null
@@ -1,140 +0,0 @@ -func @conv2d_2451x2311_same() attributes { iree.module.export } { - %inputs = iree.unfoldable_constant dense<[ - [[[ 1.0], [ 2.0], [ 3.0], [ 4.0], [ 5.0]], - [[ 6.0], [ 7.0], [ 8.0], [ 9.0], [10.0]], - [[11.0], [12.0], [13.0], [14.0], [15.0]], - [[16.0], [17.0], [18.0], [19.0], [20.0]]], - [[[21.0], [22.0], [23.0], [24.0], [25.0]], - [[26.0], [27.0], [28.0], [29.0], [30.0]], - [[31.0], [32.0], [33.0], [34.0], [35.0]], - [[36.0], [37.0], [38.0], [39.0], [40.0]]]]> : tensor <2x4x5x1xf32> - %weights = iree.unfoldable_constant dense<[ - [[[1.0]], [[2.0]], [[3.0]]], - [[[4.0]], [[5.0]], [[6.0]]]]> : tensor <2x3x1x1xf32> - %res = "mhlo.convolution"(%inputs, %weights) { - batch_group_count = 1 : i64, - dimension_numbers = { - input_batch_dimension = 0 : i64, - input_feature_dimension = 3 : i64, - input_spatial_dimensions = dense<[1, 2]> : tensor<2xi64>, - kernel_input_feature_dimension = 2 : i64, - kernel_output_feature_dimension = 3 : i64, - kernel_spatial_dimensions = dense<[0, 1]> : tensor<2xi64>, - output_batch_dimension = 0 : i64, - output_feature_dimension = 3 : i64, - output_spatial_dimensions = dense<[1, 2]> : tensor<2xi64>}, - feature_group_count = 1 : i64, - padding = dense<[[0, 1], [1, 1]]> : tensor<2x2xi64>, - rhs_dilation = dense<1> : tensor<2xi64>, - window_strides = dense<1> : tensor<2xi64>} : - (tensor<2x4x5x1xf32>, tensor<2x3x1x1xf32>) -> tensor<2x4x5x1xf32> - check.expect_almost_eq_const(%res, dense<[ - [[[ 80.0], [121.0], [142.0], [163.0], [100.0]], - [[160.0], [226.0], [247.0], [268.0], [160.0]], - [[240.0], [331.0], [352.0], [373.0], [220.0]], - [[ 83.0], [104.0], [110.0], [116.0], [ 59.0]]], - [[[400.0], [541.0], [562.0], [583.0], [340.0]], - [[480.0], [646.0], [667.0], [688.0], [400.0]], - [[560.0], [751.0], [772.0], [793.0], [460.0]], - [[183.0], [224.0], [230.0], [236.0], [119.0]]]]> : tensor<2x4x5x1xf32>) : tensor<2x4x5x1xf32> - return -} - -func @conv2d_no_padding() attributes { iree.module.export } { - %inputs = iree.unfoldable_constant dense<[ - [[[ 1.0, 2.0, 3.0], - [ 4.0, 5.0, 6.0], - [ 7.0, 8.0, 9.0], - [ 10.0, 11.0, 12.0], - [ 13.0, 14.0, 15.0]], - [[ 16.0, 17.0, 18.0], - [ 19.0, 20.0, 21.0], - [ 22.0, 23.0, 24.0], - [ 25.0, 26.0, 27.0], - [ 28.0, 29.0, 30.0]], - [[ 31.0, 32.0, 33.0], - [ 34.0, 35.0, 36.0], - [ 37.0, 38.0, 39.0], - [ 40.0, 41.0, 42.0], - [ 43.0, 44.0, 45.0]], - [[ 46.0, 47.0, 48.0], - [ 49.0, 50.0, 51.0], - [ 52.0, 53.0, 54.0], - [ 55.0, 56.0, 57.0], - [ 58.0, 59.0, 60.0]]], - [[[ 61.0, 62.0, 63.0], - [ 64.0, 65.0, 66.0], - [ 67.0, 68.0, 69.0], - [ 70.0, 71.0, 72.0], - [ 73.0, 74.0, 75.0]], - [[ 76.0, 77.0, 78.0], - [ 79.0, 80.0, 81.0], - [ 82.0, 83.0, 84.0], - [ 85.0, 86.0, 87.0], - [ 88.0, 89.0, 90.0]], - [[ 91.0, 92.0, 93.0], - [ 94.0, 95.0, 96.0], - [ 97.0, 98.0, 99.0], - [100.0, 101.0, 102.0], - [103.0, 104.0, 105.0]], - [[106.0, 107.0, 108.0], - [109.0, 110.0, 111.0], - [112.0, 113.0, 114.0], - [115.0, 116.0, 117.0], - [118.0, 119.0, 120.0]]]]> : tensor<2x4x5x3xf32> - %weights = iree.unfoldable_constant dense<[ - [[[ 1.0, 2.0, 3.0, 4.0, 5.0, 6.0], - [ 7.0, 8.0, 9.0, 10.0, 11.0, 12.0], - [ 13.0, 14.0, 15.0, 16.0, 17.0, 18.0]], - [[ 19.0, 20.0, 21.0, 22.0, 23.0, 24.0], - [ 25.0, 26.0, 27.0, 28.0, 29.0, 30.0], - [ 31.0, 32.0, 33.0, 34.0, 35.0, 36.0]], - [[ 37.0, 38.0, 39.0, 40.0, 41.0, 42.0], - [ 43.0, 44.0, 45.0, 46.0, 47.0, 48.0], - [ 49.0, 50.0, 51.0, 52.0, 53.0, 54.0]]], - [[[ 55.0, 56.0, 57.0, 58.0, 59.0, 60.0], - [ 61.0, 62.0, 63.0, 64.0, 65.0, 66.0], - [ 67.0, 68.0, 69.0, 70.0, 71.0, 72.0]], - [[ 73.0, 74.0, 75.0, 76.0, 77.0, 78.0], - [ 79.0, 80.0, 81.0, 82.0, 83.0, 84.0], - [ 85.0, 86.0, 87.0, 88.0, 89.0, 90.0]], - [[ 91.0, 92.0, 93.0, 94.0, 95.0, 96.0], - [ 97.0, 98.0, 99.0, 100.0, 101.0, 102.0], - [103.0, 104.0, 105.0, 106.0, 107.0, 108.0]]]]> : tensor<2x3x3x6xf32> - %res = "mhlo.convolution"(%inputs, %weights) { - batch_group_count = 1 : i64, - dimension_numbers = { - input_batch_dimension = 0 : i64, - input_feature_dimension = 3 : i64, - input_spatial_dimensions = dense<[1, 2]> : tensor<2xi64>, - kernel_input_feature_dimension = 2 : i64, - kernel_output_feature_dimension = 3 : i64, - kernel_spatial_dimensions = dense<[0, 1]> : tensor<2xi64>, - output_batch_dimension = 0 : i64, - output_feature_dimension = 3 : i64, - output_spatial_dimensions = dense<[1, 2]> : tensor<2xi64>}, - feature_group_count = 1 : i64, - rhs_dilation = dense<1> : tensor<2xi64>, - window_strides = dense<1> : tensor<2xi64>} : - (tensor<2x4x5x3xf32>, tensor<2x3x3x6xf32>) -> tensor<2x3x3x6xf32> - check.expect_almost_eq_const(%res, dense<[ - [[[16065.0, 16290.0, 16515.0, 16740.0, 16965.0, 17190.0], - [18873.0, 19152.0, 19431.0, 19710.0, 19989.0, 20268.0], - [21681.0, 22014.0, 22347.0, 22680.0, 23013.0, 23346.0]], - [[30105.0, 30600.0, 31095.0, 31590.0, 32085.0, 32580.0], - [32913.0, 33462.0, 34011.0, 34560.0, 35109.0, 35658.0], - [35721.0, 36324.0, 36927.0, 37530.0, 38133.0, 38736.0]], - [[44145.0, 44910.0, 45675.0, 46440.0, 47205.0, 47970.0], - [46953.0, 47772.0, 48591.0, 49410.0, 50229.0, 51048.0], - [49761.0, 50634.0, 51507.0, 52380.0, 53253.0, 54126.0]]], - [[[72225.0, 73530.0, 74835.0, 76140.0, 77445.0, 78750.0], - [75033.0, 76392.0, 77751.0, 79110.0, 80469.0, 81828.0], - [77841.0, 79254.0, 80667.0, 82080.0, 83493.0, 84906.0]], - [[86265.0, 87840.0, 89415.0, 90990.0, 92565.0, 94140.0], - [89073.0, 90702.0, 92331.0, 93960.0, 95589.0, 97218.0], - [91881.0, 93564.0, 95247.0, 96930.0, 98613.0, 100296.0]], - [[100305.0, 102150.0, 103995.0, 105840.0, 107685.0, 109530.0], - [103113.0, 105012.0, 106911.0, 108810.0, 110709.0, 112608.0], - [105921.0, 107874.0, 109827.0, 111780.0, 113733.0, 115686.0]]]]> : tensor<2x3x3x6xf32>) : tensor<2x3x3x6xf32> - return -}
diff --git a/iree/test/e2e/xla_ops/convolution.mlir b/iree/test/e2e/xla_ops/convolution.mlir index 22c5258..6ac1719 100644 --- a/iree/test/e2e/xla_ops/convolution.mlir +++ b/iree/test/e2e/xla_ops/convolution.mlir
@@ -65,51 +65,47 @@ return } -// TODO(#2345): This test seems to fail when executed with another -// test from this file, but passes as a standalone test. Needs further -// investigation - -// func @conv2d_2451x2311_same() attributes { iree.module.export } { -// %inputs = iree.unfoldable_constant dense<[ -// [[[ 1.0], [ 2.0], [ 3.0], [ 4.0], [ 5.0]], -// [[ 6.0], [ 7.0], [ 8.0], [ 9.0], [10.0]], -// [[11.0], [12.0], [13.0], [14.0], [15.0]], -// [[16.0], [17.0], [18.0], [19.0], [20.0]]], -// [[[21.0], [22.0], [23.0], [24.0], [25.0]], -// [[26.0], [27.0], [28.0], [29.0], [30.0]], -// [[31.0], [32.0], [33.0], [34.0], [35.0]], -// [[36.0], [37.0], [38.0], [39.0], [40.0]]]]> : tensor <2x4x5x1xf32> -// %weights = iree.unfoldable_constant dense<[ -// [[[1.0]], [[2.0]], [[3.0]]], -// [[[4.0]], [[5.0]], [[6.0]]]]> : tensor <2x3x1x1xf32> -// %res = "mhlo.convolution"(%inputs, %weights) { -// batch_group_count = 1 : i64, -// dimension_numbers = { -// input_batch_dimension = 0 : i64, -// input_feature_dimension = 3 : i64, -// input_spatial_dimensions = dense<[1, 2]> : tensor<2xi64>, -// kernel_input_feature_dimension = 2 : i64, -// kernel_output_feature_dimension = 3 : i64, -// kernel_spatial_dimensions = dense<[0, 1]> : tensor<2xi64>, -// output_batch_dimension = 0 : i64, -// output_feature_dimension = 3 : i64, -// output_spatial_dimensions = dense<[1, 2]> : tensor<2xi64>}, -// feature_group_count = 1 : i64, -// padding = dense<[[0, 1], [1, 1]]> : tensor<2x2xi64>, -// rhs_dilation = dense<1> : tensor<2xi64>, -// window_strides = dense<1> : tensor<2xi64>} : -// (tensor<2x4x5x1xf32>, tensor<2x3x1x1xf32>) -> tensor<2x4x5x1xf32> -// check.expect_almost_eq_const(%res, dense<[ -// [[[ 80.0], [121.0], [142.0], [163.0], [100.0]], -// [[160.0], [226.0], [247.0], [268.0], [160.0]], -// [[240.0], [331.0], [352.0], [373.0], [220.0]], -// [[ 83.0], [104.0], [110.0], [116.0], [ 59.0]]], -// [[[400.0], [541.0], [562.0], [583.0], [340.0]], -// [[480.0], [646.0], [667.0], [688.0], [400.0]], -// [[560.0], [751.0], [772.0], [793.0], [460.0]], -// [[183.0], [224.0], [230.0], [236.0], [119.0]]]]> : tensor<2x4x5x1xf32>) : tensor<2x4x5x1xf32> -// return -// } +func @conv2d_2451x2311_same() attributes { iree.module.export } { + %inputs = iree.unfoldable_constant dense<[ + [[[ 1.0], [ 2.0], [ 3.0], [ 4.0], [ 5.0]], + [[ 6.0], [ 7.0], [ 8.0], [ 9.0], [10.0]], + [[11.0], [12.0], [13.0], [14.0], [15.0]], + [[16.0], [17.0], [18.0], [19.0], [20.0]]], + [[[21.0], [22.0], [23.0], [24.0], [25.0]], + [[26.0], [27.0], [28.0], [29.0], [30.0]], + [[31.0], [32.0], [33.0], [34.0], [35.0]], + [[36.0], [37.0], [38.0], [39.0], [40.0]]]]> : tensor <2x4x5x1xf32> + %weights = iree.unfoldable_constant dense<[ + [[[1.0]], [[2.0]], [[3.0]]], + [[[4.0]], [[5.0]], [[6.0]]]]> : tensor <2x3x1x1xf32> + %res = "mhlo.convolution"(%inputs, %weights) { + batch_group_count = 1 : i64, + dimension_numbers = { + input_batch_dimension = 0 : i64, + input_feature_dimension = 3 : i64, + input_spatial_dimensions = dense<[1, 2]> : tensor<2xi64>, + kernel_input_feature_dimension = 2 : i64, + kernel_output_feature_dimension = 3 : i64, + kernel_spatial_dimensions = dense<[0, 1]> : tensor<2xi64>, + output_batch_dimension = 0 : i64, + output_feature_dimension = 3 : i64, + output_spatial_dimensions = dense<[1, 2]> : tensor<2xi64>}, + feature_group_count = 1 : i64, + padding = dense<[[0, 1], [1, 1]]> : tensor<2x2xi64>, + rhs_dilation = dense<1> : tensor<2xi64>, + window_strides = dense<1> : tensor<2xi64>} : + (tensor<2x4x5x1xf32>, tensor<2x3x1x1xf32>) -> tensor<2x4x5x1xf32> + check.expect_almost_eq_const(%res, dense<[ + [[[ 80.0], [121.0], [142.0], [163.0], [100.0]], + [[160.0], [226.0], [247.0], [268.0], [160.0]], + [[240.0], [331.0], [352.0], [373.0], [220.0]], + [[ 83.0], [104.0], [110.0], [116.0], [ 59.0]]], + [[[400.0], [541.0], [562.0], [583.0], [340.0]], + [[480.0], [646.0], [667.0], [688.0], [400.0]], + [[560.0], [751.0], [772.0], [793.0], [460.0]], + [[183.0], [224.0], [230.0], [236.0], [119.0]]]]> : tensor<2x4x5x1xf32>) : tensor<2x4x5x1xf32> + return +} func @conv2d_no_padding2() attributes { iree.module.export } { %inputs = iree.unfoldable_constant dense<[
diff --git a/iree/test/e2e/xla_ops/iota.mlir b/iree/test/e2e/xla_ops/iota.mlir new file mode 100644 index 0000000..c3ff040 --- /dev/null +++ b/iree/test/e2e/xla_ops/iota.mlir
@@ -0,0 +1,16 @@ +func @iota_dim0() attributes { iree.module.export } { + %result = "mhlo.iota"() {iota_dimension = 0 : i64} : () -> tensor<2x3xf32> + check.expect_almost_eq_const(%result, dense<[ + [0.0, 0.0, 0.0], + [1.0, 1.0, 1.0]]> : tensor<2x3xf32>) : tensor<2x3xf32> + return +} + + +func @iota_dim1() attributes { iree.module.export } { + %result = "mhlo.iota"() {iota_dimension = 1 : i64} : () -> tensor<2x3xf32> + check.expect_almost_eq_const(%result, dense<[ + [0.0, 1.0, 2.0], + [0.0, 1.0, 2.0]]> : tensor<2x3xf32>) : tensor<2x3xf32> + return +}
diff --git a/iree/test/e2e/xla_ops/pad.mlir b/iree/test/e2e/xla_ops/pad.mlir index 7f6df37..537e684 100644 --- a/iree/test/e2e/xla_ops/pad.mlir +++ b/iree/test/e2e/xla_ops/pad.mlir
@@ -20,3 +20,19 @@ check.expect_eq(%res, %input) : tensor<2x3xi32> return } + +func @pad_with_interior_padding() attributes { iree.module.export } { + %input = iree.unfoldable_constant dense<[[1, 2, 3], [4, 5, 6]]> : tensor<2x3xi32> + %c0 = iree.unfoldable_constant dense<0> : tensor<i32> + %res = "mhlo.pad"(%input, %c0) { + edge_padding_low = dense<[0, 1]> : tensor<2xi64>, + edge_padding_high = dense<[1, 5]> : tensor<2xi64>, + interior_padding = dense<[1, 2]> : tensor<2xi64> + } : (tensor<2x3xi32>, tensor<i32>) -> tensor<4x13xi32> + check.expect_eq_const(%res, dense<[ + [0, 1, 0, 0, 2, 0, 0, 3, 0, 0, 0, 0, 0], + [0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0], + [0, 4, 0, 0, 5, 0, 0, 6, 0, 0, 0, 0, 0], + [0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0]]> : tensor<4x13xi32>) : tensor<4x13xi32> + return +}
diff --git a/iree/tools/BUILD b/iree/tools/BUILD index 0f59359..46ea8a9 100644 --- a/iree/tools/BUILD +++ b/iree/tools/BUILD
@@ -50,6 +50,7 @@ "//iree/base:tracing", "//iree/modules/hal", "//iree/testing:benchmark_main", + "//iree/vm", "//iree/vm:bytecode_module", ] + PLATFORM_VULKAN_DEPS + IREE_DRIVER_MODULES, ) @@ -254,7 +255,6 @@ "//iree/modules/hal", "//iree/vm", "//iree/vm:bytecode_module", - "//iree/vm:value", "@llvm-project//llvm:Support", "@llvm-project//mlir:IR", "@llvm-project//mlir:SCFTransforms", @@ -282,6 +282,7 @@ "//iree/base:status", "//iree/base:tracing", "//iree/modules/hal", + "//iree/vm", "//iree/vm:bytecode_module", ] + PLATFORM_VULKAN_DEPS + IREE_DRIVER_MODULES, ) @@ -349,9 +350,9 @@ "//iree/base:status", "//iree/hal:api", "//iree/modules/hal", + "//iree/vm", "//iree/vm:bytecode_module", - "//iree/vm:module", - "//iree/vm:variant_list", + "//iree/vm:ref_cc", "@com_google_absl//absl/strings", "@com_google_absl//absl/types:span", ], @@ -368,7 +369,6 @@ "//iree/hal/vmla:vmla_driver_module", "//iree/modules/hal", "//iree/testing:gtest_main", - "//iree/vm:value", - "//iree/vm:variant_list", + "//iree/vm", ], )
diff --git a/iree/tools/CMakeLists.txt b/iree/tools/CMakeLists.txt index a946695..93f0b2f 100644 --- a/iree/tools/CMakeLists.txt +++ b/iree/tools/CMakeLists.txt
@@ -44,6 +44,7 @@ iree::base::tracing iree::modules::hal iree::testing::benchmark_main + iree::vm iree::vm::bytecode_module ${IREE_HAL_DRIVER_MODULES} TESTONLY @@ -84,6 +85,7 @@ iree::base::status iree::base::tracing iree::modules::hal + iree::vm iree::vm::bytecode_module ${IREE_HAL_DRIVER_MODULES} ) @@ -330,7 +332,6 @@ iree::modules::hal iree::vm iree::vm::bytecode_module - iree::vm::value ${IREE_HAL_DRIVER_MODULES} HOSTONLY ) @@ -367,9 +368,9 @@ iree::base::status iree::hal::api iree::modules::hal + iree::vm iree::vm::bytecode_module - iree::vm::module - iree::vm::variant_list + iree::vm::ref_cc PUBLIC ) @@ -386,8 +387,7 @@ iree::hal::vmla::vmla_driver_module iree::modules::hal iree::testing::gtest_main - iree::vm::value - iree::vm::variant_list + iree::vm ) if(${IREE_ENABLE_MLIR})
diff --git a/iree/tools/benchmark_module_main.cc b/iree/tools/benchmark_module_main.cc index 6f11e4a..c6cc98f 100644 --- a/iree/tools/benchmark_module_main.cc +++ b/iree/tools/benchmark_module_main.cc
@@ -22,6 +22,7 @@ #include "iree/base/tracing.h" #include "iree/modules/hal/hal_module.h" #include "iree/tools/vm_util.h" +#include "iree/vm/api.h" #include "iree/vm/bytecode_module.h" // TODO(gcmn): Allow stdin in a non-gross way. The benchmark framework invokes @@ -107,11 +108,11 @@ RETURN_IF_ERROR(ValidateFunctionAbi(function)); ASSIGN_OR_RETURN(auto input_descs, ParseInputSignature(function)); - iree_vm_variant_list_t* inputs; + vm::ref<iree_vm_list_t> inputs; if (!absl::GetFlag(FLAGS_inputs_file).empty()) { if (!absl::GetFlag(FLAGS_inputs).empty()) { return InvalidArgumentErrorBuilder(IREE_LOC) - << "Expected only one of inputs and inputs_flag to be set"; + << "Expected only one of inputs and inputs_file to be set"; } ASSIGN_OR_RETURN(inputs, ParseToVariantListFromFile( input_descs, iree_hal_device_allocator(device), @@ -124,32 +125,33 @@ ASSIGN_OR_RETURN(auto output_descs, ParseOutputSignature(function)); - iree_vm_variant_list_t* outputs = nullptr; - // Execute once to make sure any first-iteration outliers are eliminated (e.g. // JITing the SPIR-V) and clearly separate out benchmark-related problems in // future debugging. - RETURN_IF_ERROR( - FromApiStatus(iree_vm_variant_list_alloc(output_descs.size(), - IREE_ALLOCATOR_SYSTEM, &outputs), - IREE_LOC)); - RETURN_IF_ERROR( - FromApiStatus(iree_vm_invoke(context, function, /*policy=*/nullptr, - inputs, outputs, IREE_ALLOCATOR_SYSTEM), - IREE_LOC)); - iree_vm_variant_list_free(outputs); + { + vm::ref<iree_vm_list_t> outputs; + RETURN_IF_ERROR(FromApiStatus( + iree_vm_list_create(/*element_type=*/nullptr, output_descs.size(), + IREE_ALLOCATOR_SYSTEM, &outputs), + IREE_LOC)); + RETURN_IF_ERROR(FromApiStatus( + iree_vm_invoke(context, function, /*policy=*/nullptr, inputs.get(), + outputs.get(), IREE_ALLOCATOR_SYSTEM), + IREE_LOC)); + } for (auto _ : state) { // No status conversions and conditional returns in the benchmarked inner // loop. - IREE_CHECK_OK(iree_vm_variant_list_alloc(output_descs.size(), - IREE_ALLOCATOR_SYSTEM, &outputs)); - IREE_CHECK_OK(iree_vm_invoke(context, function, /*policy=*/nullptr, inputs, - outputs, IREE_ALLOCATOR_SYSTEM)); - iree_vm_variant_list_free(outputs); + vm::ref<iree_vm_list_t> outputs; + IREE_CHECK_OK(iree_vm_list_create(/*element_type=*/nullptr, + output_descs.size(), + IREE_ALLOCATOR_SYSTEM, &outputs)); + IREE_CHECK_OK(iree_vm_invoke(context, function, /*policy=*/nullptr, + inputs.get(), outputs.get(), + IREE_ALLOCATOR_SYSTEM)); } - iree_vm_variant_list_free(inputs); iree_vm_module_release(hal_module); iree_vm_module_release(input_module); iree_hal_device_release(device);
diff --git a/iree/tools/run_mlir_main.cc b/iree/tools/run_mlir_main.cc index ed09735..8a4b2c2 100644 --- a/iree/tools/run_mlir_main.cc +++ b/iree/tools/run_mlir_main.cc
@@ -66,7 +66,6 @@ #include "iree/tools/vm_util.h" #include "iree/vm/api.h" #include "iree/vm/bytecode_module.h" -#include "iree/vm/value.h" #include "llvm/ADT/StringRef.h" #include "llvm/Support/CommandLine.h" #include "llvm/Support/InitLLVM.h" @@ -274,43 +273,40 @@ std::cout << "EXEC @" << export_name << std::endl; ASSIGN_OR_RETURN(auto input_descs, ParseInputSignature(function)); - iree_vm_variant_list_t* input_list; + vm::ref<iree_vm_list_t> inputs; if (!input_values_file_flag.empty()) { if (!input_values_flag.empty()) { return InvalidArgumentErrorBuilder(IREE_LOC) - << "Expected only one of input_values_file_flag and " - "input_values_flag to be set"; + << "Expected only one of input_values and " + "input_values_file to be set"; } - ASSIGN_OR_RETURN(input_list, + ASSIGN_OR_RETURN(inputs, ParseToVariantListFromFile(input_descs, allocator, input_values_file_flag)); } else { auto input_values_list = absl::MakeConstSpan( input_values_flag.empty() ? nullptr : &input_values_flag.front(), input_values_flag.size()); - ASSIGN_OR_RETURN(input_list, ParseToVariantList(input_descs, allocator, - input_values_list)); + ASSIGN_OR_RETURN( + inputs, ParseToVariantList(input_descs, allocator, input_values_list)); } ASSIGN_OR_RETURN(auto output_descs, ParseOutputSignature(function)); // Prepare outputs list to accept the results from the invocation. - iree_vm_variant_list_t* output_list = nullptr; + vm::ref<iree_vm_list_t> outputs; RETURN_IF_ERROR(FromApiStatus( - iree_vm_variant_list_alloc(output_descs.size(), IREE_ALLOCATOR_SYSTEM, - &output_list), + iree_vm_list_create(/*element_type=*/nullptr, output_descs.size(), + IREE_ALLOCATOR_SYSTEM, &outputs), IREE_LOC)); // Synchronously invoke the function. RETURN_IF_ERROR(FromApiStatus( - iree_vm_invoke(context, function, /*policy=*/nullptr, input_list, - output_list, IREE_ALLOCATOR_SYSTEM), + iree_vm_invoke(context, function, /*policy=*/nullptr, inputs.get(), + outputs.get(), IREE_ALLOCATOR_SYSTEM), IREE_LOC)); - iree_vm_variant_list_free(input_list); - // Print outputs. - RETURN_IF_ERROR(PrintVariantList(output_descs, output_list)); - iree_vm_variant_list_free(output_list); + RETURN_IF_ERROR(PrintVariantList(output_descs, outputs.get())); return OkStatus(); }
diff --git a/iree/tools/run_module_main.cc b/iree/tools/run_module_main.cc index 45e2326..41baed1 100644 --- a/iree/tools/run_module_main.cc +++ b/iree/tools/run_module_main.cc
@@ -24,6 +24,7 @@ #include "iree/base/tracing.h" #include "iree/modules/hal/hal_module.h" #include "iree/tools/vm_util.h" +#include "iree/vm/api.h" #include "iree/vm/bytecode_module.h" ABSL_FLAG(std::string, input_file, "-", @@ -109,11 +110,11 @@ RETURN_IF_ERROR(ValidateFunctionAbi(function)); ASSIGN_OR_RETURN(auto input_descs, ParseInputSignature(function)); - iree_vm_variant_list_t* inputs; + vm::ref<iree_vm_list_t> inputs; if (!absl::GetFlag(FLAGS_inputs_file).empty()) { if (!absl::GetFlag(FLAGS_inputs).empty()) { return InvalidArgumentErrorBuilder(IREE_LOC) - << "Expected only one of inputs and inputs_flag to be set"; + << "Expected only one of inputs and inputs_file to be set"; } ASSIGN_OR_RETURN(inputs, ParseToVariantListFromFile( input_descs, iree_hal_device_allocator(device), @@ -125,24 +126,24 @@ } ASSIGN_OR_RETURN(auto output_descs, ParseOutputSignature(function)); - iree_vm_variant_list_t* outputs = nullptr; - RETURN_IF_ERROR( - FromApiStatus(iree_vm_variant_list_alloc(output_descs.size(), - IREE_ALLOCATOR_SYSTEM, &outputs), - IREE_LOC)); + vm::ref<iree_vm_list_t> outputs; + RETURN_IF_ERROR(FromApiStatus( + iree_vm_list_create(/*element_type=*/nullptr, output_descs.size(), + IREE_ALLOCATOR_SYSTEM, &outputs), + IREE_LOC)); std::cout << "EXEC @" << function_name << "\n"; - RETURN_IF_ERROR( - FromApiStatus(iree_vm_invoke(context, function, /*policy=*/nullptr, - inputs, outputs, IREE_ALLOCATOR_SYSTEM), - IREE_LOC)) + RETURN_IF_ERROR(FromApiStatus( + iree_vm_invoke(context, function, /*policy=*/nullptr, inputs.get(), + outputs.get(), IREE_ALLOCATOR_SYSTEM), + IREE_LOC)) << "invoking function " << function_name; - RETURN_IF_ERROR(PrintVariantList(output_descs, outputs)) + RETURN_IF_ERROR(PrintVariantList(output_descs, outputs.get())) << "printing results"; - iree_vm_variant_list_free(inputs); - iree_vm_variant_list_free(outputs); + inputs.reset(); + outputs.reset(); iree_vm_module_release(hal_module); iree_vm_module_release(input_module); iree_hal_device_release(device);
diff --git a/iree/tools/vm_util.cc b/iree/tools/vm_util.cc index e7c6e7a..24da00d 100644 --- a/iree/tools/vm_util.cc +++ b/iree/tools/vm_util.cc
@@ -28,8 +28,6 @@ #include "iree/hal/api.h" #include "iree/modules/hal/hal_module.h" #include "iree/vm/bytecode_module.h" -#include "iree/vm/module.h" -#include "iree/vm/variant_list.h" namespace iree { @@ -82,28 +80,19 @@ return output_descs; } -StatusOr<iree_vm_variant_list_t*> ParseToVariantListFromFile( - absl::Span<const RawSignatureParser::Description> descs, - iree_hal_allocator_t* allocator, const std::string& filename) { - ASSIGN_OR_RETURN(auto s, file_io::GetFileContents(filename)); - std::vector<std::string> input_strings = - absl::StrSplit(s, '\n', absl::SkipEmpty()); - return ParseToVariantList(descs, allocator, input_strings); -} - -StatusOr<iree_vm_variant_list_t*> ParseToVariantList( +StatusOr<vm::ref<iree_vm_list_t>> ParseToVariantList( absl::Span<const RawSignatureParser::Description> descs, iree_hal_allocator_t* allocator, - absl::Span<const std::string> input_strings) { + absl::Span<const absl::string_view> input_strings) { if (input_strings.size() != descs.size()) { return InvalidArgumentErrorBuilder(IREE_LOC) << "Signature mismatch; expected " << descs.size() << " buffer strings but received " << input_strings.size(); } - iree_vm_variant_list_t* variant_list = nullptr; + vm::ref<iree_vm_list_t> variant_list; RETURN_IF_ERROR(FromApiStatus( - iree_vm_variant_list_alloc(input_strings.size(), IREE_ALLOCATOR_SYSTEM, - &variant_list), + iree_vm_list_create(/*element_type=*/nullptr, input_strings.size(), + IREE_ALLOCATOR_SYSTEM, &variant_list), IREE_LOC)); for (int i = 0; i < input_strings.size(); ++i) { auto input_string = input_strings[i]; @@ -124,14 +113,14 @@ << "Parsing '" << input_string << "'. Has i32 descriptor but does not start with 'i32='"; } - int32_t val; - if (!absl::SimpleAtoi(input_view, &val)) { + iree_vm_value_t val = iree_vm_value_make_i32(0); + if (!absl::SimpleAtoi(input_view, &val.i32)) { return InvalidArgumentErrorBuilder(IREE_LOC) << "Converting '" << input_view << "' to i32 when parsing '" << input_string << "'"; } - iree_vm_variant_list_append_value(variant_list, - iree_vm_value_make_i32(val)); + RETURN_IF_ERROR(FromApiStatus( + iree_vm_list_push_value(variant_list.get(), &val), IREE_LOC)); break; } case RawSignatureParser::Type::kBuffer: { @@ -144,9 +133,9 @@ << "Parsing value '" << input_string << "'"; } auto buffer_view_ref = iree_hal_buffer_view_move_ref(buffer_view); - RETURN_IF_ERROR(FromApiStatus(iree_vm_variant_list_append_ref_move( - variant_list, &buffer_view_ref), - IREE_LOC)); + RETURN_IF_ERROR(FromApiStatus( + iree_vm_list_push_ref_move(variant_list.get(), &buffer_view_ref), + IREE_LOC)); break; } default: @@ -154,18 +143,36 @@ << "Unsupported signature type: " << desc_str; } } - return variant_list; + return variant_list.release(); +} + +StatusOr<vm::ref<iree_vm_list_t>> ParseToVariantList( + absl::Span<const RawSignatureParser::Description> descs, + iree_hal_allocator_t* allocator, + absl::Span<const std::string> input_strings) { + absl::InlinedVector<absl::string_view, 4> input_views(input_strings.size()); + for (int i = 0; i < input_strings.size(); ++i) { + input_views[i] = input_strings[i]; + } + return ParseToVariantList(descs, allocator, input_views); +} + +StatusOr<vm::ref<iree_vm_list_t>> ParseToVariantListFromFile( + absl::Span<const RawSignatureParser::Description> descs, + iree_hal_allocator_t* allocator, const std::string& filename) { + ASSIGN_OR_RETURN(auto file_string, file_io::GetFileContents(filename)); + absl::InlinedVector<absl::string_view, 4> input_views( + absl::StrSplit(file_string, '\n', absl::SkipEmpty())); + return ParseToVariantList(descs, allocator, input_views); } Status PrintVariantList(absl::Span<const RawSignatureParser::Description> descs, - iree_vm_variant_list_t* variant_list, - std::ostream* os) { - for (int i = 0; i < iree_vm_variant_list_size(variant_list); ++i) { - iree_vm_variant_t* variant = iree_vm_variant_list_get(variant_list, i); - if (!variant) { - return InvalidArgumentErrorBuilder(IREE_LOC) - << "variant " << i << "not present"; - } + iree_vm_list_t* variant_list, std::ostream* os) { + for (int i = 0; i < iree_vm_list_size(variant_list); ++i) { + iree_vm_variant_t variant = iree_vm_variant_empty(); + RETURN_IF_ERROR(FromApiStatus( + iree_vm_list_get_variant(variant_list, i, &variant), IREE_LOC)) + << "variant " << i << "not present"; const auto& desc = descs[i]; std::string desc_str; @@ -174,27 +181,27 @@ switch (desc.type) { case RawSignatureParser::Type::kScalar: { - if (variant->value_type != IREE_VM_VALUE_TYPE_I32) { + if (variant.type.value_type != IREE_VM_VALUE_TYPE_I32) { return InvalidArgumentErrorBuilder(IREE_LOC) << "variant " << i << " has value type " - << static_cast<int>(variant->value_type) + << static_cast<int>(variant.type.value_type) << " but descriptor information " << desc_str; } if (desc.scalar.type != AbiConstants::ScalarType::kSint32) { return UnimplementedErrorBuilder(IREE_LOC) << "Unsupported signature scalar type: " << desc_str; } - *os << "i32=" << variant->i32 << "\n"; + *os << "i32=" << variant.i32 << "\n"; break; } case RawSignatureParser::Type::kBuffer: { - if (variant->value_type != IREE_VM_VALUE_TYPE_NONE) { + if (!iree_vm_type_def_is_ref(&variant.type)) { return InvalidArgumentErrorBuilder(IREE_LOC) << "variant " << i << " has value type " - << static_cast<int>(variant->value_type) + << static_cast<int>(variant.type.value_type) << " but descriptor information " << desc_str; } - auto* buffer_view = iree_hal_buffer_view_deref(&variant->ref); + auto* buffer_view = iree_hal_buffer_view_deref(&variant.ref); if (!buffer_view) { return InvalidArgumentErrorBuilder(IREE_LOC) << "failed dereferencing variant " << i;
diff --git a/iree/tools/vm_util.h b/iree/tools/vm_util.h index c3f1c9c..ac8dbac 100644 --- a/iree/tools/vm_util.h +++ b/iree/tools/vm_util.h
@@ -22,8 +22,8 @@ #include "iree/base/signature_mangle.h" #include "iree/base/status.h" #include "iree/hal/api.h" -#include "iree/vm/module.h" -#include "iree/vm/variant_list.h" +#include "iree/vm/api.h" +#include "iree/vm/ref_cc.h" namespace iree { @@ -40,13 +40,6 @@ StatusOr<std::vector<RawSignatureParser::Description>> ParseOutputSignature( const iree_vm_function_t& function); -// Parses the content in |filename| into a variant list of VM scalars and -// buffers. See ParseToVariantList for the format of scalars and buffers. The -// inputs are expected to be newline-separated. -StatusOr<iree_vm_variant_list_t*> ParseToVariantListFromFile( - absl::Span<const RawSignatureParser::Description> descs, - iree_hal_allocator_t* allocator, const std::string& filename); - // Parses |input_strings| into a variant list of VM scalars and buffers. // Scalars should be in the format: // type=value @@ -57,11 +50,22 @@ // Uses |allocator| to allocate the buffers. // Uses descriptors in |descs| for type information and validation. // The returned variant list must be freed by the caller. -StatusOr<iree_vm_variant_list_t*> ParseToVariantList( +StatusOr<vm::ref<iree_vm_list_t>> ParseToVariantList( + absl::Span<const RawSignatureParser::Description> descs, + iree_hal_allocator_t* allocator, + absl::Span<const absl::string_view> input_strings); +StatusOr<vm::ref<iree_vm_list_t>> ParseToVariantList( absl::Span<const RawSignatureParser::Description> descs, iree_hal_allocator_t* allocator, absl::Span<const std::string> input_strings); +// Parses the content in |filename| into a variant list of VM scalars and +// buffers. See ParseToVariantList for the format of scalars and buffers. The +// inputs are expected to be newline-separated. +StatusOr<vm::ref<iree_vm_list_t>> ParseToVariantListFromFile( + absl::Span<const RawSignatureParser::Description> descs, + iree_hal_allocator_t* allocator, const std::string& filename); + // Prints a variant list of VM scalars and buffers to |os|. // Prints scalars in the format: // type=value @@ -71,7 +75,7 @@ // https://github.com/google/iree/tree/main/iree/hal/api.h // Uses descriptors in |descs| for type information and validation. Status PrintVariantList(absl::Span<const RawSignatureParser::Description> descs, - iree_vm_variant_list_t* variant_list, + iree_vm_list_t* variant_list, std::ostream* os = &std::cout); // Creates the default device for |driver| in |out_device|.
diff --git a/iree/tools/vm_util_test.cc b/iree/tools/vm_util_test.cc index 11a0092..2391fbb 100644 --- a/iree/tools/vm_util_test.cc +++ b/iree/tools/vm_util_test.cc
@@ -21,8 +21,7 @@ #include "iree/hal/api.h" #include "iree/modules/hal/hal_module.h" #include "iree/testing/gtest.h" -#include "iree/vm/value.h" -#include "iree/vm/variant_list.h" +#include "iree/vm/api.h" namespace iree { namespace { @@ -42,72 +41,64 @@ }; TEST_F(VmUtilTest, ParsePrintBuffer) { - auto buf_string = "2x2xi32=[42 43][44 45]"; + absl::string_view buf_string = "2x2xi32=[42 43][44 45]"; RawSignatureParser::Description desc; desc.type = RawSignatureParser::Type::kBuffer; desc.buffer.scalar_type = AbiConstants::ScalarType::kSint32; desc.dims = {2, 2}; - ASSERT_OK_AND_ASSIGN(auto* variant_list, + ASSERT_OK_AND_ASSIGN(auto variant_list, ParseToVariantList({desc}, allocator_, {buf_string})); std::stringstream os; - ASSERT_OK(PrintVariantList({desc}, variant_list, &os)); + ASSERT_OK(PrintVariantList({desc}, variant_list.get(), &os)); EXPECT_EQ(os.str(), absl::StrCat(buf_string, "\n")); - - iree_vm_variant_list_free(variant_list); } TEST_F(VmUtilTest, ParsePrintScalar) { - auto input_string = "i32=42"; + absl::string_view input_string = "i32=42"; RawSignatureParser::Description desc; desc.type = RawSignatureParser::Type::kScalar; desc.scalar.type = AbiConstants::ScalarType::kSint32; - ASSERT_OK_AND_ASSIGN(auto* variant_list, + ASSERT_OK_AND_ASSIGN(auto variant_list, ParseToVariantList({desc}, allocator_, {input_string})); std::stringstream os; - ASSERT_OK(PrintVariantList({desc}, variant_list, &os)); + ASSERT_OK(PrintVariantList({desc}, variant_list.get(), &os)); EXPECT_EQ(os.str(), absl::StrCat(input_string, "\n")); - - iree_vm_variant_list_free(variant_list); } TEST_F(VmUtilTest, ParsePrintRank0Buffer) { - auto buf_string = "i32=42"; + absl::string_view buf_string = "i32=42"; RawSignatureParser::Description desc; desc.type = RawSignatureParser::Type::kBuffer; desc.buffer.scalar_type = AbiConstants::ScalarType::kSint32; - ASSERT_OK_AND_ASSIGN(auto* variant_list, + ASSERT_OK_AND_ASSIGN(auto variant_list, ParseToVariantList({desc}, allocator_, {buf_string})); std::stringstream os; - ASSERT_OK(PrintVariantList({desc}, variant_list, &os)); + ASSERT_OK(PrintVariantList({desc}, variant_list.get(), &os)); EXPECT_EQ(os.str(), absl::StrCat(buf_string, "\n")); - - iree_vm_variant_list_free(variant_list); } TEST_F(VmUtilTest, ParsePrintMultipleBuffers) { - auto buf_string1 = "2x2xi32=[42 43][44 45]"; + absl::string_view buf_string1 = "2x2xi32=[42 43][44 45]"; RawSignatureParser::Description desc1; desc1.type = RawSignatureParser::Type::kBuffer; desc1.buffer.scalar_type = AbiConstants::ScalarType::kSint32; desc1.dims = {2, 2}; - auto buf_string2 = "2x3xf64=[1 2 3][4 5 6]"; + absl::string_view buf_string2 = "2x3xf64=[1 2 3][4 5 6]"; RawSignatureParser::Description desc2; desc2.type = RawSignatureParser::Type::kBuffer; desc2.buffer.scalar_type = AbiConstants::ScalarType::kIeeeFloat64; desc2.dims = {2, 3}; - ASSERT_OK_AND_ASSIGN(auto* variant_list, + ASSERT_OK_AND_ASSIGN(auto variant_list, ParseToVariantList({desc1, desc2}, allocator_, {buf_string1, buf_string2})); std::stringstream os; - ASSERT_OK(PrintVariantList({desc1, desc2}, variant_list, &os)); + ASSERT_OK(PrintVariantList({desc1, desc2}, variant_list.get(), &os)); EXPECT_EQ(os.str(), absl::StrCat(buf_string1, "\n", buf_string2, "\n")); - - iree_vm_variant_list_free(variant_list); } } // namespace
diff --git a/iree/vm/BUILD b/iree/vm/BUILD index c9350c5..93829d0 100644 --- a/iree/vm/BUILD +++ b/iree/vm/BUILD
@@ -92,6 +92,24 @@ ) cc_test( + name = "bytecode_module_size_benchmark", + srcs = ["bytecode_module_size_benchmark.cc"], + deps = [ + ":bytecode_module", + ":bytecode_module_size_benchmark_module_cc", + ":vm", + "//iree/base:api", + ], +) + +iree_bytecode_module( + name = "bytecode_module_size_benchmark_module", + src = "bytecode_module_size_benchmark.mlir", + cc_namespace = "iree::vm", + flags = ["-iree-vm-ir-to-bytecode-module"], +) + +cc_test( name = "bytecode_module_test", srcs = ["bytecode_module_test.cc"], deps = [ @@ -149,8 +167,8 @@ hdrs = ["invocation.h"], deps = [ ":context", + ":list", ":module", - ":variant_list", "//iree/base:api", "//iree/base:tracing", ], @@ -251,7 +269,6 @@ deps = [ ":module", ":ref", - ":variant_list", "//iree/base:alignment", "//iree/base:api", "//iree/base:tracing", @@ -285,16 +302,6 @@ ) cc_library( - name = "variant_list", - srcs = ["variant_list.c"], - hdrs = ["variant_list.h"], - deps = [ - ":ref", - ":value", - ], -) - -cc_library( name = "vm", hdrs = [ "api.h", @@ -310,7 +317,6 @@ ":stack", ":type_def", ":value", - ":variant_list", "//iree/base:api", ], )
diff --git a/iree/vm/CMakeLists.txt b/iree/vm/CMakeLists.txt index 45c96c6..ae91187 100644 --- a/iree/vm/CMakeLists.txt +++ b/iree/vm/CMakeLists.txt
@@ -106,6 +106,30 @@ iree_cc_test( NAME + bytecode_module_size_benchmark + SRCS + "bytecode_module_size_benchmark.cc" + DEPS + ::bytecode_module + ::bytecode_module_size_benchmark_module_cc + ::vm + iree::base::api +) + +iree_bytecode_module( + NAME + bytecode_module_size_benchmark_module + SRC + "bytecode_module_size_benchmark.mlir" + CC_NAMESPACE + "iree::vm" + FLAGS + "-iree-vm-ir-to-bytecode-module" + PUBLIC +) + +iree_cc_test( + NAME bytecode_module_test SRCS "bytecode_module_test.cc" @@ -165,8 +189,8 @@ "invocation.c" DEPS ::context + ::list ::module - ::variant_list iree::base::api iree::base::tracing PUBLIC @@ -283,7 +307,6 @@ DEPS ::module ::ref - ::variant_list iree::base::alignment iree::base::api iree::base::tracing @@ -324,19 +347,6 @@ iree_cc_library( NAME - variant_list - HDRS - "variant_list.h" - SRCS - "variant_list.c" - DEPS - ::ref - ::value - PUBLIC -) - -iree_cc_library( - NAME vm HDRS "api.h" @@ -351,7 +361,6 @@ ::stack ::type_def ::value - ::variant_list iree::base::api PUBLIC )
diff --git a/iree/vm/api.h b/iree/vm/api.h index 57331ae..e4f2f4e 100644 --- a/iree/vm/api.h +++ b/iree/vm/api.h
@@ -20,10 +20,11 @@ #include "iree/vm/context.h" #include "iree/vm/instance.h" #include "iree/vm/invocation.h" +#include "iree/vm/list.h" #include "iree/vm/module.h" #include "iree/vm/ref.h" #include "iree/vm/stack.h" +#include "iree/vm/type_def.h" #include "iree/vm/value.h" -#include "iree/vm/variant_list.h" #endif // IREE_VM_API_H_
diff --git a/iree/vm/bytecode_dispatch.c b/iree/vm/bytecode_dispatch.c index 01e5b2d..eab5136 100644 --- a/iree/vm/bytecode_dispatch.c +++ b/iree/vm/bytecode_dispatch.c
@@ -63,6 +63,13 @@ static_assert(offsetof(iree_vm_register_remap_list_t, pairs) == 2, "Expect no padding in the struct"); +// Maps a type ID to a type def with clamping for out of bounds values. +static inline const iree_vm_type_def_t* iree_vm_map_type( + iree_vm_bytecode_module_t* module, int32_t type_id) { + type_id = type_id >= module->type_count ? 0 : type_id; + return &module->type_table[type_id]; +} + // Remaps registers from a source set to a destination set within the same stack // frame. This is a way to perform a conditional multi-mov sequence instead of // requiring the additional bytecode representation of the conditional movs. @@ -103,6 +110,76 @@ } } +static const int kRegSize = sizeof(uint16_t); + +// Bytecode data access macros for reading values of a given type from a byte +// offset within the current function. +#if defined(IREE_IS_LITTLE_ENDIAN) +#define OP_I8(i) bytecode_data[pc + i] +#define OP_I16(i) *((uint16_t*)&bytecode_data[pc + i]) +#define OP_I32(i) *((uint32_t*)&bytecode_data[pc + i]) +#else +#define OP_I8(i) bytecode_data[pc + i] +#define OP_I16(i) \ + ((uint16_t)bytecode_data[pc + 0 + i]) | \ + ((uint16_t)bytecode_data[pc + 1 + i] << 8) +#define OP_I32(i) \ + ((uint32_t)bytecode_data[pc + 0 + i]) | \ + ((uint32_t)bytecode_data[pc + 1 + i] << 8) | \ + ((uint32_t)bytecode_data[pc + 2 + i] << 16) | \ + ((uint32_t)bytecode_data[pc + 3 + i] << 24) +#endif // IREE_IS_LITTLE_ENDIAN + +// These utilities match the VM_Enc* statements in VMBase.td 1:1, allowing us +// to have the inverse of the encoding which make things easier to read. +// +// Each macro will increment the pc by the number of bytes read and as such must +// be called in the same order the values are encoded. +#define VM_DecConstI8(name) \ + OP_I8(0); \ + ++pc; +#define VM_DecConstI32(name) \ + OP_I32(0); \ + pc += 4; +#define VM_DecOpcode(opcode) VM_DecConstI8(#opcode) +#define VM_DecFuncAttr(name) VM_DecConstI32(name) +#define VM_DecGlobalAttr(name) VM_DecConstI32(name) +#define VM_DecRodataAttr(name) VM_DecConstI32(name) +#define VM_DecType(name) \ + iree_vm_map_type(module, OP_I32(0)); \ + pc += 4; +#define VM_DecTypeOf(name) VM_DecType(name) +#define VM_DecIntAttr32(name) VM_DecConstI32(name) +#define VM_DecStrAttr(name, out_str) \ + (out_str)->size = (iree_host_size_t)OP_I16(0); \ + (out_str)->data = (const char*)&bytecode_data[pc + 2]; \ + pc += 2 + (out_str)->size; +#define VM_DecBranchTarget(block_name) VM_DecConstI32(name) +#define VM_DecBranchOperands(operands_name) \ + (const iree_vm_register_remap_list_t*)&bytecode_data[pc]; \ + pc += \ + kRegSize + ((const iree_vm_register_list_t*)&bytecode_data[pc])->size * \ + 2 * kRegSize; +#define VM_DecOperandRegI32(name) \ + regs.i32[OP_I16(0) & regs.i32_mask]; \ + pc += kRegSize; +#define VM_DecOperandRegRef(name, out_is_move) \ + ®s.ref[OP_I16(0) & regs.ref_mask]; \ + *(out_is_move) = OP_I16(0) & IREE_REF_REGISTER_MOVE_BIT; \ + pc += kRegSize; +#define VM_DecVariadicOperands(name) \ + (const iree_vm_register_list_t*)&bytecode_data[pc]; \ + pc += kRegSize + \ + ((const iree_vm_register_list_t*)&bytecode_data[pc])->size * kRegSize; +#define VM_DecResultRegI32(name) \ + ®s.i32[OP_I16(0) & regs.i32_mask]; \ + pc += kRegSize; +#define VM_DecResultRegRef(name, out_is_move) \ + ®s.ref[OP_I16(0) & regs.ref_mask]; \ + *(out_is_move) = OP_I16(0) & IREE_REF_REGISTER_MOVE_BIT; \ + pc += kRegSize; +#define VM_DecVariadicResults(name) VM_DecVariadicOperands(name) + iree_status_t iree_vm_bytecode_dispatch( iree_vm_bytecode_module_t* module, iree_vm_bytecode_module_state_t* module_state, iree_vm_stack_t* stack, @@ -132,7 +209,7 @@ #define DECLARE_DISPATCH_OPC(ordinal, name) &&_dispatch_##name, #define DECLARE_DISPATCH_RSV(ordinal) &&_dispatch_unhandled, static const void* kDispatchTable[256] = { - IREE_VM_OP_TABLE(DECLARE_DISPATCH_OPC, DECLARE_DISPATCH_RSV)}; + IREE_VM_OP_CORE_TABLE(DECLARE_DISPATCH_OPC, DECLARE_DISPATCH_RSV)}; #define DISPATCH_UNHANDLED() \ _dispatch_unhandled: \ @@ -161,35 +238,13 @@ return IREE_STATUS_UNIMPLEMENTED; #define DISPATCH_OP(op_name, body) \ - case IREE_VM_OP_##op_name: \ + case IREE_VM_OP_CORE_##op_name: \ IREE_DISPATCH_LOG_OPCODE(#op_name); \ body; \ break; #endif // IREE_DISPATCH_MODE_COMPUTED_GOTO - static const int kRegSize = sizeof(uint16_t); - -#if defined(IREE_IS_LITTLE_ENDIAN) -#define OP_I8(i) bytecode_data[pc + i] -#define OP_I16(i) *((uint16_t*)&bytecode_data[pc + i]) -#define OP_I32(i) *((uint32_t*)&bytecode_data[pc + i]) -#else -#define OP_I8(i) bytecode_data[pc + i] -#define OP_I16(i) \ - ((uint16_t)bytecode_data[pc + 0 + i]) | \ - ((uint16_t)bytecode_data[pc + 1 + i] << 8) -#define OP_I32(i) \ - ((uint32_t)bytecode_data[pc + 0 + i]) | \ - ((uint32_t)bytecode_data[pc + 1 + i] << 8) | \ - ((uint32_t)bytecode_data[pc + 2 + i] << 16) | \ - ((uint32_t)bytecode_data[pc + 3 + i] << 24) -#endif // IREE_IS_LITTLE_ENDIAN - -#define OP_R_I32(i) regs.i32[OP_I16(i) & regs.i32_mask] -#define OP_R_REF(i) regs.ref[OP_I16(i) & regs.ref_mask] -#define OP_R_REF_IS_MOVE(i) (OP_I16(i) & IREE_REF_REGISTER_MOVE_BIT) - // Primary dispatch state. This is our 'native stack frame' and really // just enough to make dereferencing common addresses (like the current // offset) faster. You can think of this like CPU state (like PC). @@ -208,159 +263,109 @@ memset(out_result, 0, sizeof(*out_result)); - // NOTE: we should generate this with tblgen, as it has the encoding info. - // TODO(benvanik): at least generate operand reading/writing and sizes. - // This could look something like: - // OP_GlobalLoadI32_value = OP_GlobalLoadI32_global; - // pc += OP_Size_GlobalLoadI32; - BEGIN_DISPATCH() { //===------------------------------------------------------------------===// // Globals //===------------------------------------------------------------------===// DISPATCH_OP(GlobalLoadI32, { - // let encoding = [ - // VM_EncOpcode<VM_OPC_GlobalLoadI32>, - // VM_EncGlobalAttr<"global">, - // VM_EncResult<"value">, - // ]; - int byte_offset = OP_I32(0); + int32_t byte_offset = VM_DecGlobalAttr("global"); if (byte_offset < 0 || byte_offset >= module_state->rwdata_storage.data_length) { return IREE_STATUS_OUT_OF_RANGE; } + int32_t* value = VM_DecResultRegI32("value"); const int32_t* global_ptr = (const int32_t*)(module_state->rwdata_storage.data + byte_offset); - OP_R_I32(4) = *global_ptr; - pc += 4 + kRegSize; + *value = *global_ptr; }); + DISPATCH_OP(GlobalStoreI32, { - // let encoding = [ - // VM_EncOpcode<VM_OPC_GlobalStoreI32>, - // VM_EncGlobalAttr<"global">, - // VM_EncOperand<"value", 0>, - // ]; - int byte_offset = OP_I32(0); + int32_t byte_offset = VM_DecGlobalAttr("global"); if (byte_offset < 0 || byte_offset >= module_state->rwdata_storage.data_length) { return IREE_STATUS_OUT_OF_RANGE; } + int32_t value = VM_DecOperandRegI32("value"); int32_t* global_ptr = (int32_t*)(module_state->rwdata_storage.data + byte_offset); - *global_ptr = OP_R_I32(4); - pc += 4 + kRegSize; + *global_ptr = value; }); DISPATCH_OP(GlobalLoadIndirectI32, { - // let encoding = [ - // VM_EncOpcode<VM_OPC_GlobalLoadIndirectI32>, - // VM_EncOperand<"global", 0>, - // VM_EncResult<"value">, - // ]; - int byte_offset = OP_R_I32(0); + int32_t byte_offset = VM_DecOperandRegI32("global"); if (byte_offset < 0 || byte_offset >= module_state->rwdata_storage.data_length) { return IREE_STATUS_OUT_OF_RANGE; } + int32_t* value = VM_DecResultRegI32("value"); const int32_t* global_ptr = (const int32_t*)(module_state->rwdata_storage.data + byte_offset); - OP_R_I32(2) = *global_ptr; - pc += kRegSize + kRegSize; + *value = *global_ptr; }); + DISPATCH_OP(GlobalStoreIndirectI32, { - // let encoding = [ - // VM_EncOpcode<VM_OPC_GlobalStoreIndirectI32>, - // VM_EncOperand<"global", 0>, - // VM_EncOperand<"value", 1>, - // ]; - int byte_offset = OP_R_I32(0); + int32_t byte_offset = VM_DecOperandRegI32("global"); if (byte_offset < 0 || byte_offset >= module_state->rwdata_storage.data_length) { return IREE_STATUS_OUT_OF_RANGE; } + int32_t value = VM_DecOperandRegI32("value"); int32_t* global_ptr = (int32_t*)(module_state->rwdata_storage.data + byte_offset); - *global_ptr = OP_R_I32(2); - pc += kRegSize + kRegSize; + *global_ptr = value; }); DISPATCH_OP(GlobalLoadRef, { - // let encoding = [ - // VM_EncOpcode<VM_OPC_GlobalLoadRef>, - // VM_EncGlobalAttr<"global">, - // VM_EncTypeOf<"value">, - // VM_EncResult<"value">, - // ]; - int global = OP_I32(0); + int32_t global = VM_DecGlobalAttr("global"); if (global < 0 || global >= module_state->global_ref_count) { return IREE_STATUS_OUT_OF_RANGE; } - int type_id = OP_I32(4); - type_id = type_id >= module->type_count ? 0 : type_id; - const iree_vm_type_def_t* type_def = &module->type_table[type_id]; + const iree_vm_type_def_t* type_def = VM_DecTypeOf("value"); + bool result_is_move; + iree_vm_ref_t* result = VM_DecResultRegRef("value", &result_is_move); iree_vm_ref_t* global_ref = &module_state->global_ref_table[global]; - iree_vm_ref_retain_or_move_checked(OP_R_REF_IS_MOVE(8), global_ref, - type_def->ref_type, &OP_R_REF(8)); - pc += 4 + 4 + kRegSize; + iree_vm_ref_retain_or_move_checked(result_is_move, global_ref, + type_def->ref_type, result); }); + DISPATCH_OP(GlobalStoreRef, { - // let encoding = [ - // VM_EncOpcode<VM_OPC_GlobalStoreRef>, - // VM_EncGlobalAttr<"global">, - // VM_EncTypeOf<"value">, - // VM_EncOperand<"value", 0>, - // ]; - int global = OP_I32(0); + int32_t global = VM_DecGlobalAttr("global"); if (global < 0 || global >= module_state->global_ref_count) { return IREE_STATUS_OUT_OF_RANGE; } - int type_id = OP_I32(4); - type_id = type_id >= module->type_count ? 0 : type_id; - const iree_vm_type_def_t* type_def = &module->type_table[type_id]; + const iree_vm_type_def_t* type_def = VM_DecTypeOf("value"); + bool value_is_move; + iree_vm_ref_t* value = VM_DecOperandRegRef("value", &value_is_move); iree_vm_ref_t* global_ref = &module_state->global_ref_table[global]; - iree_vm_ref_retain_or_move_checked(OP_R_REF_IS_MOVE(8), &OP_R_REF(8), + iree_vm_ref_retain_or_move_checked(value_is_move, value, type_def->ref_type, global_ref); - pc += 4 + 4 + kRegSize; }); DISPATCH_OP(GlobalLoadIndirectRef, { - // let encoding = [ - // VM_EncOpcode<VM_OPC_GlobalLoadIndirectRef>, - // VM_EncOperand<"global", 0>, - // VM_EncTypeOf<"value">, - // VM_EncResult<"value">, - // ]; - int global = OP_R_I32(0); + int32_t global = VM_DecGlobalAttr("global"); if (global < 0 || global >= module_state->global_ref_count) { return IREE_STATUS_OUT_OF_RANGE; } - int type_id = OP_I32(2); - type_id = type_id >= module->type_count ? 0 : type_id; - const iree_vm_type_def_t* type_def = &module->type_table[type_id]; + const iree_vm_type_def_t* type_def = VM_DecTypeOf("value"); + bool result_is_move; + iree_vm_ref_t* result = VM_DecResultRegRef("value", &result_is_move); iree_vm_ref_t* global_ref = &module_state->global_ref_table[global]; - iree_vm_ref_retain_or_move_checked(OP_R_REF_IS_MOVE(6), global_ref, - type_def->ref_type, &OP_R_REF(6)); - pc += kRegSize + 4 + kRegSize; + iree_vm_ref_retain_or_move_checked(result_is_move, global_ref, + type_def->ref_type, result); }); + DISPATCH_OP(GlobalStoreIndirectRef, { - // let encoding = [ - // VM_EncOpcode<VM_OPC_GlobalStoreIndirectRef>, - // VM_EncOperand<"global", 0>, - // VM_EncTypeOf<"value">, - // VM_EncOperand<"value", 1>, - // ]; - int global = OP_R_I32(0); + int32_t global = VM_DecGlobalAttr("global"); if (global < 0 || global >= module_state->global_ref_count) { return IREE_STATUS_OUT_OF_RANGE; } - int type_id = OP_I32(2); - type_id = type_id >= module->type_count ? 0 : type_id; - const iree_vm_type_def_t* type_def = &module->type_table[type_id]; + const iree_vm_type_def_t* type_def = VM_DecTypeOf("value"); + bool value_is_move; + iree_vm_ref_t* value = VM_DecOperandRegRef("value", &value_is_move); iree_vm_ref_t* global_ref = &module_state->global_ref_table[global]; - iree_vm_ref_retain_or_move_checked(OP_R_REF_IS_MOVE(6), &OP_R_REF(6), + iree_vm_ref_retain_or_move_checked(value_is_move, value, type_def->ref_type, global_ref); - pc += kRegSize + 4 + kRegSize; }); //===------------------------------------------------------------------===// @@ -368,48 +373,32 @@ //===------------------------------------------------------------------===// DISPATCH_OP(ConstI32, { - // let encoding = [ - // VM_EncOpcode<opcode>, - // VM_EncIntAttr<"value", type.bitwidth>, - // VM_EncResult<"result">, - // ]; - OP_R_I32(4) = OP_I32(0); - pc += 4 + kRegSize; + int32_t value = VM_DecIntAttr32("value"); + int32_t* result = VM_DecResultRegI32("result"); + *result = value; }); DISPATCH_OP(ConstI32Zero, { - // let encoding = [ - // VM_EncOpcode<VM_OPC_ConstI32Zero>, - // VM_EncResult<"result">, - // ]; - OP_R_I32(0) = 0; - pc += kRegSize; + int32_t* result = VM_DecResultRegI32("result"); + *result = 0; }); DISPATCH_OP(ConstRefZero, { - // let encoding = [ - // VM_EncOpcode<VM_OPC_ConstRefZero>, - // VM_EncResult<"result">, - // ]; - iree_vm_ref_release(&OP_R_REF(0)); - pc += kRegSize; + bool result_is_move; + iree_vm_ref_t* result = VM_DecResultRegRef("result", &result_is_move); + iree_vm_ref_release(result); }); DISPATCH_OP(ConstRefRodata, { - // let encoding = [ - // VM_EncOpcode<VM_OPC_ConstRefRodata>, - // VM_EncRodataAttr<"rodata">, - // VM_EncResult<"value">, - // ]; - int32_t rodata_ordinal = OP_I32(0); + int32_t rodata_ordinal = VM_DecRodataAttr("rodata"); if (rodata_ordinal < 0 || rodata_ordinal >= module_state->rodata_ref_count) { return IREE_STATUS_OUT_OF_RANGE; } - // TODO(benvanik): allow decompression callbacks to run now (if needed). + bool result_is_move; + iree_vm_ref_t* result = VM_DecResultRegRef("value", &result_is_move); iree_vm_ref_wrap_retain(&module_state->rodata_ref_table[rodata_ordinal], - iree_vm_ro_byte_buffer_type_id(), &OP_R_REF(4)); - pc += 4 + kRegSize; + iree_vm_ro_byte_buffer_type_id(), result); }); //===------------------------------------------------------------------===// @@ -417,120 +406,86 @@ //===------------------------------------------------------------------===// DISPATCH_OP(ListAlloc, { - // let encoding = [ - // VM_EncOpcode<VM_OPC_ListAlloc>, - // VM_EncType<"{{element_type}}">, - // VM_EncOperand<"initial_capacity", 0>, - // VM_EncResult<"result">, - // ]; - int element_type_id = OP_I32(0); - element_type_id = - element_type_id >= module->type_count ? 0 : element_type_id; - const iree_vm_type_def_t* element_type_def = - &module->type_table[element_type_id]; - iree_host_size_t initial_capacity = (iree_host_size_t)OP_R_I32(4); + const iree_vm_type_def_t* element_type_def = VM_DecTypeOf("element_type"); + iree_host_size_t initial_capacity = + VM_DecOperandRegI32("initial_capacity"); + bool result_is_move; + iree_vm_ref_t* result = VM_DecResultRegRef("result", &result_is_move); iree_vm_list_t* list = NULL; IREE_RETURN_IF_ERROR(iree_vm_list_create( element_type_def, initial_capacity, module_state->allocator, &list)); - iree_vm_ref_wrap_assign(list, iree_vm_list_type_id(), &OP_R_REF(6)); - pc += 4 + kRegSize + kRegSize; + iree_vm_ref_wrap_assign(list, iree_vm_list_type_id(), result); }); DISPATCH_OP(ListReserve, { - // let encoding = [ - // VM_EncOpcode<VM_OPC_ListReserve>, - // VM_EncOperand<"list", 0>, - // VM_EncOperand<"minimum_capacity", 1>, - // ]; - iree_vm_ref_t* list_ref = &OP_R_REF(0); + bool list_is_move; + iree_vm_ref_t* list_ref = VM_DecOperandRegRef("list", &list_is_move); iree_vm_list_t* list = iree_vm_list_deref(list_ref); if (!list) return IREE_STATUS_INVALID_ARGUMENT; - IREE_RETURN_IF_ERROR(iree_vm_list_reserve(list, OP_R_I32(2))); - pc += kRegSize + kRegSize; + int32_t minimum_capacity = VM_DecOperandRegI32("minimum_capacity"); + IREE_RETURN_IF_ERROR(iree_vm_list_reserve(list, minimum_capacity)); }); DISPATCH_OP(ListSize, { - // let encoding = [ - // VM_EncOpcode<VM_OPC_ListSize>, - // VM_EncOperand<"list", 0>, - // VM_EncResult<"result">, - // ]; - iree_vm_ref_t* list_ref = &OP_R_REF(0); + bool list_is_move; + iree_vm_ref_t* list_ref = VM_DecOperandRegRef("list", &list_is_move); iree_vm_list_t* list = iree_vm_list_deref(list_ref); if (!list) return IREE_STATUS_INVALID_ARGUMENT; - OP_R_I32(2) = (int32_t)iree_vm_list_size(list); - pc += kRegSize + kRegSize; + int32_t* result = VM_DecResultRegI32("result"); + *result = (int32_t)iree_vm_list_size(list); }); DISPATCH_OP(ListResize, { - // let encoding = [ - // VM_EncOpcode<VM_OPC_ListResize>, - // VM_EncOperand<"list", 0>, - // VM_EncOperand<"new_size", 1>, - // ]; - iree_vm_ref_t* list_ref = &OP_R_REF(0); + bool list_is_move; + iree_vm_ref_t* list_ref = VM_DecOperandRegRef("list", &list_is_move); iree_vm_list_t* list = iree_vm_list_deref(list_ref); if (!list) return IREE_STATUS_INVALID_ARGUMENT; - IREE_RETURN_IF_ERROR(iree_vm_list_resize(list, OP_R_I32(2))); - pc += kRegSize + kRegSize; + int32_t new_size = VM_DecOperandRegI32("new_size"); + IREE_RETURN_IF_ERROR(iree_vm_list_resize(list, new_size)); }); DISPATCH_OP(ListGetI32, { - // let encoding = [ - // VM_EncOpcode<opcode>, - // VM_EncOperand<"list", 0>, - // VM_EncOperand<"index", 1>, - // VM_EncResult<"result">, - // ]; - iree_vm_ref_t* list_ref = &OP_R_REF(0); + bool list_is_move; + iree_vm_ref_t* list_ref = VM_DecOperandRegRef("list", &list_is_move); iree_vm_list_t* list = iree_vm_list_deref(list_ref); if (!list) return IREE_STATUS_INVALID_ARGUMENT; + int32_t index = VM_DecOperandRegI32("index"); + int32_t* result = VM_DecResultRegI32("result"); iree_vm_value_t value; IREE_RETURN_IF_ERROR(iree_vm_list_get_value_as( - list, OP_R_I32(2), IREE_VM_VALUE_TYPE_I32, &value)); - OP_R_I32(4) = value.i32; - pc += kRegSize + kRegSize + kRegSize; + list, index, IREE_VM_VALUE_TYPE_I32, &value)); + *result = value.i32; }); DISPATCH_OP(ListSetI32, { - // let encoding = [ - // VM_EncOpcode<opcode>, - // VM_EncOperand<"list", 0>, - // VM_EncOperand<"index", 1>, - // VM_EncOperand<"value", 2>, - // ]; - iree_vm_ref_t* list_ref = &OP_R_REF(0); + bool list_is_move; + iree_vm_ref_t* list_ref = VM_DecOperandRegRef("list", &list_is_move); iree_vm_list_t* list = iree_vm_list_deref(list_ref); if (!list) return IREE_STATUS_INVALID_ARGUMENT; - iree_vm_value_t value = iree_vm_value_make_i32(OP_R_I32(4)); - IREE_RETURN_IF_ERROR(iree_vm_list_set_value(list, OP_R_I32(2), &value)); - pc += kRegSize + kRegSize + kRegSize; + int32_t index = VM_DecOperandRegI32("index"); + int32_t raw_value = VM_DecOperandRegI32("raw_value"); + iree_vm_value_t value = iree_vm_value_make_i32(raw_value); + IREE_RETURN_IF_ERROR(iree_vm_list_set_value(list, index, &value)); }); DISPATCH_OP(ListGetRef, { - // let encoding = [ - // VM_EncOpcode<VM_OPC_ListGetRef>, - // VM_EncOperand<"list", 0>, - // VM_EncOperand<"index", 1>, - // VM_EncTypeOf<"result">, - // VM_EncResult<"result">, - // ]; - iree_vm_ref_t* list_ref = &OP_R_REF(0); - iree_vm_list_t* list = iree_vm_list_deref(list_ref); - if (!list) return IREE_STATUS_INVALID_ARGUMENT; + // bool list_is_move; + // iree_vm_ref_t* list_ref = VM_DecOperandRegRef("list", &list_is_move); + // iree_vm_list_t* list = iree_vm_list_deref(list_ref); + // if (!list) return IREE_STATUS_INVALID_ARGUMENT; + // int32_t index = VM_DecOperandRegI32("index"); + // iree_vm_ref_t* result = VM_DecResultRegRef("result"); return IREE_STATUS_UNIMPLEMENTED; }); DISPATCH_OP(ListSetRef, { - // let encoding = [ - // VM_EncOpcode<VM_OPC_ListSetRef>, - // VM_EncOperand<"list", 0>, - // VM_EncOperand<"index", 1>, - // VM_EncOperand<"value", 2>, - // ]; - iree_vm_ref_t* list_ref = &OP_R_REF(0); - iree_vm_list_t* list = iree_vm_list_deref(list_ref); - if (!list) return IREE_STATUS_INVALID_ARGUMENT; + // bool list_is_move; + // iree_vm_ref_t* list_ref = VM_DecOperandRegRef("list", &list_is_move); + // iree_vm_list_t* list = iree_vm_list_deref(list_ref); + // if (!list) return IREE_STATUS_INVALID_ARGUMENT; + // int32_t index = VM_DecOperandRegI32("index"); + // bool operand_is_move = VM_DecOperandRegRefIsMove("value"); + // iree_vm_ref_t* operand = VM_DecOperandRegRef("value"); return IREE_STATUS_UNIMPLEMENTED; }); @@ -539,122 +494,92 @@ //===------------------------------------------------------------------===// DISPATCH_OP(SelectI32, { - // let encoding = [ - // VM_EncOpcode<opcode>, - // VM_EncOperand<"condition", 0>, - // VM_EncOperand<"true_value", 1>, - // VM_EncOperand<"false_value", 2>, - // VM_EncResult<"result">, - // ]; - OP_R_I32(6) = OP_R_I32(0) ? OP_R_I32(2) : OP_R_I32(4); - pc += kRegSize + kRegSize + kRegSize + kRegSize; + int32_t condition = VM_DecOperandRegI32("condition"); + int32_t true_value = VM_DecOperandRegI32("true_value"); + int32_t false_value = VM_DecOperandRegI32("false_value"); + int32_t* result = VM_DecResultRegI32("result"); + *result = condition ? true_value : false_value; }); DISPATCH_OP(SelectRef, { - // let encoding = [ - // VM_EncOpcode<VM_OPC_SelectRef>, - // VM_EncOperand<"condition", 0>, - // VM_EncTypeOf<"true_value">, - // VM_EncOperand<"true_value", 1>, - // VM_EncOperand<"false_value", 2>, - // VM_EncResult<"result">, - // ]; + int32_t condition = VM_DecOperandRegI32("condition"); // TODO(benvanik): remove the type_id and use either LHS/RHS (if both are // null then output is always null so no need to know the type). - int type_id = OP_I32(2); - type_id = type_id >= module->type_count ? 0 : type_id; - const iree_vm_type_def_t* type_def = &module->type_table[type_id]; - if (OP_R_I32(0)) { - // Select LHS (+6). - iree_vm_ref_retain_or_move_checked(OP_R_REF_IS_MOVE(6), &OP_R_REF(6), - type_def->ref_type, &OP_R_REF(10)); - if (OP_R_REF_IS_MOVE(10)) iree_vm_ref_release(&OP_R_REF(8)); + const iree_vm_type_def_t* type_def = VM_DecTypeOf("true_value"); + bool true_value_is_move; + iree_vm_ref_t* true_value = + VM_DecOperandRegRef("true_value", &true_value_is_move); + bool false_value_is_move; + iree_vm_ref_t* false_value = + VM_DecOperandRegRef("false_value", &false_value_is_move); + bool result_is_move; + iree_vm_ref_t* result = VM_DecResultRegRef("result", &result_is_move); + if (condition) { + // Select LHS. + iree_vm_ref_retain_or_move_checked(true_value_is_move, true_value, + type_def->ref_type, result); + if (false_value_is_move) iree_vm_ref_release(false_value); } else { - // Select RHS (+8). - iree_vm_ref_retain_or_move_checked(OP_R_REF_IS_MOVE(8), &OP_R_REF(8), - type_def->ref_type, &OP_R_REF(10)); - if (OP_R_REF_IS_MOVE(6)) iree_vm_ref_release(&OP_R_REF(6)); + // Select RHS. + iree_vm_ref_retain_or_move_checked(false_value_is_move, false_value, + type_def->ref_type, result); + if (true_value_is_move) iree_vm_ref_release(true_value); } - pc += kRegSize + 4 + kRegSize + kRegSize + kRegSize; }); DISPATCH_OP(SwitchI32, { - // let encoding = [ - // VM_EncOpcode<opcode>, - // VM_EncOperand<"index", 0>, - // VM_EncIntAttr<"default_value", 32>, - // VM_EncVariadicOperands<"values">, - // VM_EncResult<"result">, - // ]; - int32_t index = OP_R_I32(0); - int32_t default_value = OP_I32(2); + int32_t index = VM_DecOperandRegI32("index"); + int32_t default_value = VM_DecIntAttr32("default_value"); const iree_vm_register_list_t* value_reg_list = - (const iree_vm_register_list_t*)&bytecode_data[pc + 6]; - pc += kRegSize + 4 + kRegSize + value_reg_list->size * kRegSize; - int32_t new_value = default_value; + VM_DecVariadicOperands("values"); + int32_t* result = VM_DecResultRegI32("result"); if (index >= 0 && index < value_reg_list->size) { - new_value = regs.i32[value_reg_list->registers[index] & regs.i32_mask]; + *result = regs.i32[value_reg_list->registers[index] & regs.i32_mask]; + } else { + *result = default_value; } - OP_R_I32(0) = new_value; - pc += kRegSize; }); DISPATCH_OP(SwitchRef, { - // let encoding = [ - // VM_EncOpcode<VM_OPC_SwitchRef>, - // VM_EncOperand<"index", 0>, - // VM_EncTypeOf<"result">, - // VM_EncOperand<"default_value", 1>, - // VM_EncVariadicOperands<"values">, - // VM_EncResult<"result">, - // ]; - int32_t index = OP_R_I32(0); - int32_t type_id = OP_I32(2); - iree_vm_ref_t* default_value = &OP_R_REF(6); - int is_move = OP_R_REF_IS_MOVE(6); + int32_t index = VM_DecOperandRegI32("index"); + const iree_vm_type_def_t* type_def = VM_DecTypeOf("result"); + bool default_is_move; + iree_vm_ref_t* default_value = + VM_DecOperandRegRef("default_value", &default_is_move); const iree_vm_register_list_t* value_reg_list = - (const iree_vm_register_list_t*)&bytecode_data[pc + 8]; - // Skip over all operands. - pc += - kRegSize + 4 + kRegSize + kRegSize + value_reg_list->size * kRegSize; - iree_vm_ref_t* new_value = default_value; + VM_DecVariadicOperands("values"); + bool result_is_move; + iree_vm_ref_t* result = VM_DecResultRegRef("result", &result_is_move); if (index >= 0 && index < value_reg_list->size) { - new_value = ®s.ref[value_reg_list->registers[index] & regs.ref_mask]; - is_move = value_reg_list->registers[index] & IREE_REF_REGISTER_MOVE_BIT; + bool is_move = + value_reg_list->registers[index] & IREE_REF_REGISTER_MOVE_BIT; + iree_vm_ref_t* new_value = + ®s.ref[value_reg_list->registers[index] & regs.ref_mask]; + iree_vm_ref_retain_or_move_checked(is_move, new_value, + type_def->ref_type, result); + } else { + iree_vm_ref_retain_or_move_checked(default_is_move, default_value, + type_def->ref_type, result); } - iree_vm_ref_t* result_reg = &OP_R_REF(0); - pc += kRegSize; - type_id = type_id >= module->type_count ? 0 : type_id; - const iree_vm_type_def_t* type_def = &module->type_table[type_id]; - iree_vm_ref_retain_or_move_checked(is_move, new_value, type_def->ref_type, - result_reg); }); //===------------------------------------------------------------------===// // Native integer arithmetic //===------------------------------------------------------------------===// - // let encoding = [ - // VM_EncOpcode<opcode>, - // VM_EncOperand<"operand", 0>, - // VM_EncResult<"result">, - // ]; -#define DISPATCH_OP_UNARY_ALU_I32(op_name, type, op) \ - DISPATCH_OP(op_name, { \ - OP_R_I32(2) = (int32_t)(op((type)OP_R_I32(0))); \ - pc += kRegSize + kRegSize; \ +#define DISPATCH_OP_UNARY_ALU_I32(op_name, type, op) \ + DISPATCH_OP(op_name, { \ + int32_t operand = VM_DecOperandRegI32("operand"); \ + int32_t* result = VM_DecResultRegI32("result"); \ + *result = (int32_t)(op((type)operand)); \ }); - // let encoding = [ - // VM_EncOpcode<opcode>, - // VM_EncOperand<"lhs", 0>, - // VM_EncOperand<"rhs", 1>, - // VM_EncResult<"result">, - // ]; -#define DISPATCH_OP_BINARY_ALU_I32(op_name, type, op) \ - DISPATCH_OP(op_name, { \ - OP_R_I32(4) = (int32_t)(((type)OP_R_I32(0))op((type)OP_R_I32(2))); \ - pc += kRegSize + kRegSize + kRegSize; \ +#define DISPATCH_OP_BINARY_ALU_I32(op_name, type, op) \ + DISPATCH_OP(op_name, { \ + int32_t lhs = VM_DecOperandRegI32("lhs"); \ + int32_t rhs = VM_DecOperandRegI32("rhs"); \ + int32_t* result = VM_DecResultRegI32("result"); \ + *result = (int32_t)(((type)lhs)op((type)rhs)); \ }); DISPATCH_OP_BINARY_ALU_I32(AddI32, int32_t, +); @@ -673,36 +598,30 @@ // Casting and type conversion/emulation //===------------------------------------------------------------------===// - // let encoding = [ - // VM_EncOpcode<opcode>, - // VM_EncOperand<"operand", 0>, - // VM_EncResult<"result">, - // ]; #define DISPATCH_OP_CAST_I32(op_name, src_type, dst_type) \ DISPATCH_OP(op_name, { \ - OP_R_I32(2) = (dst_type)((src_type)OP_R_I32(0)); \ - pc += kRegSize + kRegSize; \ + int32_t operand = VM_DecOperandRegI32("operand"); \ + int32_t* result = VM_DecResultRegI32("result"); \ + *result = (dst_type)((src_type)operand); \ }); - DISPATCH_OP_CAST_I32(TruncI8, uint8_t, uint32_t); - DISPATCH_OP_CAST_I32(TruncI16, uint16_t, uint32_t); + DISPATCH_OP_CAST_I32(TruncI32I8, uint8_t, uint32_t); + DISPATCH_OP_CAST_I32(TruncI32I16, uint16_t, uint32_t); DISPATCH_OP_CAST_I32(ExtI8I32S, int8_t, int32_t); + DISPATCH_OP_CAST_I32(ExtI8I32U, uint8_t, uint32_t); DISPATCH_OP_CAST_I32(ExtI16I32S, int16_t, int32_t); + DISPATCH_OP_CAST_I32(ExtI16I32U, uint16_t, uint32_t); //===------------------------------------------------------------------===// // Native bitwise shifts and rotates //===------------------------------------------------------------------===// - // let encoding = [ - // VM_EncOpcode<opcode>, - // VM_EncOperand<"operand", 0>, - // VM_EncIntAttr<"amount", type.bitwidth>, - // VM_EncResult<"result">, - // ]; -#define DISPATCH_OP_SHIFT_I32(op_name, type, op) \ - DISPATCH_OP(op_name, { \ - OP_R_I32(4) = (int32_t)(((type)OP_R_I32(0))op OP_I8(2)); \ - pc += kRegSize + kRegSize + kRegSize; \ +#define DISPATCH_OP_SHIFT_I32(op_name, type, op) \ + DISPATCH_OP(op_name, { \ + int32_t operand = VM_DecOperandRegI32("operand"); \ + int32_t amount = VM_DecConstI8("amount"); \ + int32_t* result = VM_DecResultRegI32("result"); \ + *result = (int32_t)(((type)operand)op amount); \ }); DISPATCH_OP_SHIFT_I32(ShlI32, int32_t, <<); @@ -713,67 +632,50 @@ // Comparison ops //===------------------------------------------------------------------===// - // let encoding = [ - // VM_EncOpcode<opcode>, - // VM_EncOperand<"lhs", 0>, - // VM_EncOperand<"rhs", 1>, - // VM_EncResult<"result">, - // ]; -#define DISPATCH_OP_CMP_I32(op_name, type, op) \ - DISPATCH_OP(op_name, { \ - OP_R_I32(4) = (((type)OP_R_I32(0))op((type)OP_R_I32(2))) ? 1 : 0; \ - pc += kRegSize + kRegSize + kRegSize; \ +#define DISPATCH_OP_CMP_I32(op_name, type, op) \ + DISPATCH_OP(op_name, { \ + int32_t lhs = VM_DecOperandRegI32("lhs"); \ + int32_t rhs = VM_DecOperandRegI32("rhs"); \ + int32_t* result = VM_DecResultRegI32("result"); \ + *result = (((type)lhs)op((type)rhs)) ? 1 : 0; \ }); DISPATCH_OP_CMP_I32(CmpEQI32, int32_t, ==); DISPATCH_OP_CMP_I32(CmpNEI32, int32_t, !=); DISPATCH_OP_CMP_I32(CmpLTI32S, int32_t, <); DISPATCH_OP_CMP_I32(CmpLTI32U, uint32_t, <); - DISPATCH_OP_CMP_I32(CmpLTEI32S, int32_t, <=); - DISPATCH_OP_CMP_I32(CmpLTEI32U, uint32_t, <=); - DISPATCH_OP_CMP_I32(CmpGTI32S, int32_t, >); - DISPATCH_OP_CMP_I32(CmpGTI32U, uint32_t, >); - DISPATCH_OP_CMP_I32(CmpGTEI32S, int32_t, >=); - DISPATCH_OP_CMP_I32(CmpGTEI32U, uint32_t, >=); DISPATCH_OP(CmpNZI32, { - OP_R_I32(2) = (OP_R_I32(0) != 0) ? 1 : 0; - pc += kRegSize + kRegSize; + int32_t operand = VM_DecOperandRegI32("operand"); + int32_t* result = VM_DecResultRegI32("result"); + *result = (operand != 0) ? 1 : 0; }); DISPATCH_OP(CmpEQRef, { - // let encoding = [ - // VM_EncOpcode<opcode>, - // VM_EncOperand<"lhs", 0>, - // VM_EncOperand<"rhs", 1>, - // VM_EncResult<"result">, - // ]; - // TODO(benvanik): move refs. - OP_R_I32(4) = iree_vm_ref_equal(&OP_R_REF(0), &OP_R_REF(2)); - if (OP_R_REF_IS_MOVE(0)) iree_vm_ref_release(&OP_R_REF(0)); - if (OP_R_REF_IS_MOVE(2)) iree_vm_ref_release(&OP_R_REF(2)); - pc += kRegSize + kRegSize + kRegSize; + bool lhs_is_move; + iree_vm_ref_t* lhs = VM_DecOperandRegRef("lhs", &lhs_is_move); + bool rhs_is_move; + iree_vm_ref_t* rhs = VM_DecOperandRegRef("rhs", &rhs_is_move); + int32_t* result = VM_DecResultRegI32("result"); + *result = iree_vm_ref_equal(lhs, rhs); + if (lhs_is_move) iree_vm_ref_release(lhs); + if (rhs_is_move) iree_vm_ref_release(rhs); }); DISPATCH_OP(CmpNERef, { - // let encoding = [ - // VM_EncOpcode<opcode>, - // VM_EncOperand<"lhs", 0>, - // VM_EncOperand<"rhs", 1>, - // VM_EncResult<"result">, - // ]; - OP_R_I32(4) = !iree_vm_ref_equal(&OP_R_REF(0), &OP_R_REF(2)); - if (OP_R_REF_IS_MOVE(0)) iree_vm_ref_release(&OP_R_REF(0)); - if (OP_R_REF_IS_MOVE(2)) iree_vm_ref_release(&OP_R_REF(2)); - pc += kRegSize + kRegSize + kRegSize; + bool lhs_is_move; + iree_vm_ref_t* lhs = VM_DecOperandRegRef("lhs", &lhs_is_move); + bool rhs_is_move; + iree_vm_ref_t* rhs = VM_DecOperandRegRef("rhs", &rhs_is_move); + int32_t* result = VM_DecResultRegI32("result"); + *result = !iree_vm_ref_equal(lhs, rhs); + if (lhs_is_move) iree_vm_ref_release(lhs); + if (rhs_is_move) iree_vm_ref_release(rhs); }); DISPATCH_OP(CmpNZRef, { - // let encoding = [ - // VM_EncOpcode<opcode>, - // VM_EncOperand<"operand", 0>, - // VM_EncResult<"result">, - // ]; - OP_R_I32(2) = OP_R_REF(0).ptr != NULL; - if (OP_R_REF_IS_MOVE(0)) iree_vm_ref_release(&OP_R_REF(0)); - pc += kRegSize + kRegSize; + bool operand_is_move; + iree_vm_ref_t* operand = VM_DecOperandRegRef("operand", &operand_is_move); + int32_t* result = VM_DecResultRegI32("result"); + *result = operand->ptr != NULL ? 1 : 0; + if (operand_is_move) iree_vm_ref_release(operand); }); //===------------------------------------------------------------------===// @@ -781,39 +683,22 @@ //===------------------------------------------------------------------===// DISPATCH_OP(Branch, { - // let encoding = [ - // VM_EncOpcode<VM_OPC_Branch>, - // VM_EncBranch<"dest", "operands">, - // ]; - - int32_t block_pc = OP_I32(0); + int32_t block_pc = VM_DecBranchTarget("dest"); const iree_vm_register_remap_list_t* remap_list = - (const iree_vm_register_remap_list_t*)&bytecode_data[pc + 4]; - pc += 4 + kRegSize + remap_list->size * 2 * kRegSize; + VM_DecBranchOperands("operands"); pc = block_pc; iree_vm_bytecode_dispatch_remap_branch_registers(regs, remap_list); }); DISPATCH_OP(CondBranch, { - // let encoding = [ - // VM_EncOpcode<VM_OPC_CondBranch>, - // VM_EncOperand<"condition", 0>, - // VM_EncBranch<"getTrueDest", "getTrueOperands">, - // VM_EncBranch<"getFalseDest", "getFalseOperands">, - // ]; - - int32_t cond_value = OP_R_I32(0); - int32_t true_block_pc = OP_I32(2); + int32_t condition = VM_DecOperandRegI32("condition"); + int32_t true_block_pc = VM_DecBranchTarget("true_dest"); const iree_vm_register_remap_list_t* true_remap_list = - (const iree_vm_register_remap_list_t*)&bytecode_data[pc + kRegSize + - 4]; - pc += kRegSize + 4 + kRegSize + true_remap_list->size * 2 * kRegSize; - int32_t false_block_pc = OP_I32(0); + VM_DecBranchOperands("true_operands"); + int32_t false_block_pc = VM_DecBranchTarget("false_dest"); const iree_vm_register_remap_list_t* false_remap_list = - (const iree_vm_register_remap_list_t*)&bytecode_data[pc + 4]; - pc += 4 + kRegSize + false_remap_list->size * 2 * kRegSize; - - if (cond_value) { + VM_DecBranchOperands("false_operands"); + if (condition) { pc = true_block_pc; iree_vm_bytecode_dispatch_remap_branch_registers(regs, true_remap_list); } else { @@ -824,21 +709,11 @@ }); DISPATCH_OP(Call, { - // let encoding = [ - // VM_EncOpcode<VM_OPC_Call>, - // VM_EncFuncAttr<"callee">, - // VM_EncVariadicOperands<"operands">, - // VM_EncVariadicResults<"results">, - // ]; - - // Get argument and result register lists and flush the caller frame. - int32_t function_ordinal = OP_I32(0); + int32_t function_ordinal = VM_DecFuncAttr("callee"); const iree_vm_register_list_t* src_reg_list = - (const iree_vm_register_list_t*)&bytecode_data[pc + 4]; - pc += 4 + kRegSize + src_reg_list->size * kRegSize; + VM_DecVariadicOperands("operands"); const iree_vm_register_list_t* dst_reg_list = - (const iree_vm_register_list_t*)&bytecode_data[pc]; - pc += kRegSize + dst_reg_list->size * kRegSize; + VM_DecVariadicResults("results"); current_frame->pc = pc; // NOTE: we assume validation has ensured these functions exist. @@ -887,29 +762,15 @@ }); DISPATCH_OP(CallVariadic, { - // let encoding = [ - // VM_EncOpcode<VM_OPC_CallVariadic>, - // VM_EncFuncAttr<"callee">, - // VM_EncIntArrayAttr<"segment_sizes", 16>, - // VM_EncVariadicOperands<"operands">, - // VM_EncVariadicResults<"results">, - // ]; - // TODO(benvanik): dedupe with above or merge and always have the seg size // list be present (but empty) for non-variadic calls. - - // Get argument and result register lists and flush the caller frame. - int32_t function_ordinal = OP_I32(0); - pc += 4; + int32_t function_ordinal = VM_DecFuncAttr("callee"); const iree_vm_register_list_t* seg_size_list = - (const iree_vm_register_list_t*)&bytecode_data[pc]; - pc += kRegSize + seg_size_list->size * kRegSize; + VM_DecVariadicOperands("segment_sizes"); const iree_vm_register_list_t* src_reg_list = - (const iree_vm_register_list_t*)&bytecode_data[pc]; - pc += kRegSize + src_reg_list->size * kRegSize; + VM_DecVariadicOperands("operands"); const iree_vm_register_list_t* dst_reg_list = - (const iree_vm_register_list_t*)&bytecode_data[pc]; - pc += kRegSize + dst_reg_list->size * kRegSize; + VM_DecVariadicResults("results"); current_frame->pc = pc; // NOTE: we assume validation has ensured these functions exist. @@ -940,15 +801,9 @@ }); DISPATCH_OP(Return, { - // let encoding = [ - // VM_EncOpcode<VM_OPC_Return>, - // VM_EncVariadicOperands<"operands">, - // ]; - - // Remap registers from callee to caller. const iree_vm_register_list_t* src_reg_list = - (const iree_vm_register_list_t*)&bytecode_data[pc]; - current_frame->pc = pc + kRegSize + src_reg_list->size * kRegSize; + VM_DecVariadicOperands("operands"); + current_frame->pc = pc; // Leave callee by cleaning up the stack. iree_vm_stack_function_leave(stack, src_reg_list, ¤t_frame); @@ -969,16 +824,9 @@ }); DISPATCH_OP(Fail, { - // let encoding = [ - // VM_EncOpcode<VM_OPC_Fail>, - // VM_EncOperand<"status", 0>, - // VM_EncStrAttr<"message">, - // ]; - uint32_t status_code = OP_R_I32(0); - iree_string_view_t str; - str.size = OP_I16(2); - str.data = (const char*)&bytecode_data[pc + 2 + 2]; - pc += 2 + 2 + str.size; + uint32_t status_code = VM_DecOperandRegI32("status"); + iree_string_view_t message; + VM_DecStrAttr("message", &message); // TODO(benvanik): attach string and stack. if (status_code == 0) { // Shouldn't happen; we expect to die here, so there's no way to no-op. @@ -992,9 +840,6 @@ //===------------------------------------------------------------------===// DISPATCH_OP(Yield, { - // let encoding = [ - // VM_EncOpcode<VM_OPC_Yield>, - // ]; // TODO(benvanik): yield with execution results. return IREE_STATUS_OK; }); @@ -1004,71 +849,54 @@ //===------------------------------------------------------------------===// DISPATCH_OP(Trace, { - // let encoding = [ - // VM_EncOpcode<VM_OPC_Trace>, - // VM_EncStrAttr<"event_name">, - // VM_EncVariadicOperands<"operands">, - // ]; - iree_string_view_t str; - str.size = OP_I16(0); - str.data = (const char*)&bytecode_data[pc + 2]; - pc += 2 + str.size; + iree_string_view_t event_name; + VM_DecStrAttr("event_name", &event_name); const iree_vm_register_list_t* src_reg_list = - (const iree_vm_register_list_t*)&bytecode_data[pc]; - pc += kRegSize + src_reg_list->size * kRegSize; + VM_DecVariadicOperands("operands"); // TODO(benvanik): trace (if enabled). iree_vm_bytecode_dispatch_discard_registers(regs, src_reg_list); }); DISPATCH_OP(Print, { - // let encoding = [ - // VM_EncOpcode<VM_OPC_Print>, - // VM_EncStrAttr<"message">, - // VM_EncVariadicOperands<"operands">, - // ]; - iree_string_view_t str; - str.size = OP_I16(0); - str.data = (const char*)&bytecode_data[pc + 2]; - pc += 2 + str.size; + iree_string_view_t event_name; + VM_DecStrAttr("event_name", &event_name); const iree_vm_register_list_t* src_reg_list = - (const iree_vm_register_list_t*)&bytecode_data[pc]; - pc += kRegSize + src_reg_list->size * kRegSize; + VM_DecVariadicOperands("operands"); // TODO(benvanik): print. iree_vm_bytecode_dispatch_discard_registers(regs, src_reg_list); }); DISPATCH_OP(Break, { - // let encoding = [ - // VM_EncOpcode<VM_OPC_Break>, - // VM_EncBranch<"dest", "operands">, - // ]; // TODO(benvanik): break unconditionally. - int32_t block_pc = OP_I32(0); + int32_t block_pc = VM_DecBranchTarget("dest"); const iree_vm_register_remap_list_t* remap_list = - (const iree_vm_register_remap_list_t*)&bytecode_data[pc + 4]; - pc += 4 + kRegSize + remap_list->size * 2 * kRegSize; + VM_DecBranchOperands("operands"); iree_vm_bytecode_dispatch_remap_branch_registers(regs, remap_list); pc = block_pc; }); DISPATCH_OP(CondBreak, { - // let encoding = [ - // VM_EncOpcode<VM_OPC_CondBreak>, - // VM_EncBranch<"dest", "operands">, - // ]; - int32_t cond_value = OP_R_I32(0); - if (cond_value) { + int32_t condition = VM_DecOperandRegI32("condition"); + if (condition) { // TODO(benvanik): cond break. } - int32_t block_pc = OP_I32(2); + int32_t block_pc = VM_DecBranchTarget("dest"); const iree_vm_register_remap_list_t* remap_list = - (const iree_vm_register_remap_list_t*)&bytecode_data[pc + kRegSize + - 4]; - pc += kRegSize + 4 + kRegSize + remap_list->size * 2 * kRegSize; + VM_DecBranchOperands("operands"); iree_vm_bytecode_dispatch_remap_branch_registers(regs, remap_list); pc = block_pc; }); + //===------------------------------------------------------------------===// + // Extension trampolines + //===------------------------------------------------------------------===// + + DISPATCH_OP(PrefixExtI64, { return IREE_STATUS_UNIMPLEMENTED; }); + + DISPATCH_OP(PrefixExtF32, { return IREE_STATUS_UNIMPLEMENTED; }); + + DISPATCH_OP(PrefixExtF64, { return IREE_STATUS_UNIMPLEMENTED; }); + // NOLINTNEXTLINE(misc-static-assert) DISPATCH_UNHANDLED(); }
diff --git a/iree/vm/bytecode_module_size_benchmark.cc b/iree/vm/bytecode_module_size_benchmark.cc new file mode 100644 index 0000000..969ef22 --- /dev/null +++ b/iree/vm/bytecode_module_size_benchmark.cc
@@ -0,0 +1,50 @@ +// Copyright 2020 Google LLC +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// https://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#include "iree/base/api.h" +#include "iree/vm/api.h" +#include "iree/vm/bytecode_module.h" +#include "iree/vm/bytecode_module_size_benchmark_module.h" + +extern "C" int main(int argc, char** argv) { + iree_vm_instance_t* instance = nullptr; + iree_vm_instance_create(IREE_ALLOCATOR_SYSTEM, &instance); + + const auto* module_file_toc = + iree::vm::bytecode_module_size_benchmark_module_create(); + iree_vm_module_t* module = nullptr; + iree_vm_bytecode_module_create( + iree_const_byte_span_t{ + reinterpret_cast<const uint8_t*>(module_file_toc->data), + module_file_toc->size}, + IREE_ALLOCATOR_NULL, IREE_ALLOCATOR_SYSTEM, &module); + + iree_vm_context_t* context = nullptr; + iree_vm_context_create_with_modules(instance, &module, /*module_count=*/1, + IREE_ALLOCATOR_SYSTEM, &context); + + iree_vm_function_t function; + iree_vm_module_lookup_function_by_name( + module, IREE_VM_FUNCTION_LINKAGE_EXPORT, + iree_make_cstring_view("empty_func"), &function); + + iree_vm_invoke(context, function, /*policy=*/nullptr, /*inputs=*/nullptr, + /*outputs=*/nullptr, IREE_ALLOCATOR_SYSTEM); + + iree_vm_module_release(module); + iree_vm_context_release(context); + iree_vm_instance_release(instance); + + return 0; +}
diff --git a/iree/vm/bytecode_module_size_benchmark.mlir b/iree/vm/bytecode_module_size_benchmark.mlir new file mode 100644 index 0000000..84ad13c --- /dev/null +++ b/iree/vm/bytecode_module_size_benchmark.mlir
@@ -0,0 +1,6 @@ +vm.module @bytecode_module_size_benchmark { + vm.export @empty_func + vm.func @empty_func() { + vm.return + } +}
diff --git a/iree/vm/invocation.c b/iree/vm/invocation.c index e04d796..f8cf3f9 100644 --- a/iree/vm/invocation.c +++ b/iree/vm/invocation.c
@@ -22,45 +22,45 @@ // Marshals a variant list of values into callee registers. // The |out_dst_reg_list| will be populated with the register ordinals and must -// be preallocated to store iree_vm_variant_list_size inputs. -static void iree_vm_stack_frame_marshal_inputs( - iree_vm_variant_list_t* inputs, const iree_vm_registers_t dst_regs, +// be preallocated to store iree_vm_list_size inputs. +static iree_status_t iree_vm_stack_frame_marshal_inputs( + iree_vm_list_t* inputs, const iree_vm_registers_t dst_regs, iree_vm_register_list_t* out_dst_reg_list) { - iree_host_size_t count = iree_vm_variant_list_size(inputs); + iree_host_size_t count = iree_vm_list_size(inputs); uint16_t i32_reg = 0; uint16_t ref_reg = 0; out_dst_reg_list->size = (uint16_t)count; for (iree_host_size_t i = 0; i < count; ++i) { - iree_vm_variant_t* variant = iree_vm_variant_list_get(inputs, i); - if (IREE_VM_VARIANT_IS_REF(variant)) { + iree_vm_variant_t variant = iree_vm_variant_empty(); + IREE_RETURN_IF_ERROR(iree_vm_list_get_variant(inputs, i, &variant)); + if (iree_vm_type_def_is_ref(&variant.type)) { out_dst_reg_list->registers[i] = ref_reg | IREE_REF_REGISTER_TYPE_BIT | IREE_REF_REGISTER_MOVE_BIT; iree_vm_ref_t* reg_ref = &dst_regs.ref[ref_reg++]; memset(reg_ref, 0, sizeof(*reg_ref)); - iree_vm_ref_retain(&variant->ref, reg_ref); + iree_vm_ref_retain(&variant.ref, reg_ref); } else { out_dst_reg_list->registers[i] = i32_reg; - dst_regs.i32[i32_reg++] = variant->i32; + dst_regs.i32[i32_reg++] = variant.i32; } } + return IREE_STATUS_OK; } // Marshals callee return registers into a variant list. static iree_status_t iree_vm_stack_frame_marshal_outputs( const iree_vm_registers_t src_regs, - const iree_vm_register_list_t* src_reg_list, - iree_vm_variant_list_t* outputs) { + const iree_vm_register_list_t* src_reg_list, iree_vm_list_t* outputs) { for (int i = 0; i < src_reg_list->size; ++i) { uint16_t reg = src_reg_list->registers[i]; if (reg & IREE_REF_REGISTER_TYPE_BIT) { iree_vm_ref_t* value = &src_regs.ref[reg & src_regs.ref_mask]; - IREE_RETURN_IF_ERROR( - iree_vm_variant_list_append_ref_move(outputs, value)); + IREE_RETURN_IF_ERROR(iree_vm_list_push_ref_move(outputs, value)); } else { iree_vm_value_t value; value.type = IREE_VM_VALUE_TYPE_I32; value.i32 = src_regs.i32[reg & src_regs.i32_mask]; - IREE_RETURN_IF_ERROR(iree_vm_variant_list_append_value(outputs, value)); + IREE_RETURN_IF_ERROR(iree_vm_list_push_value(outputs, &value)); } } return IREE_STATUS_OK; @@ -70,13 +70,12 @@ static iree_status_t iree_vm_invoke_within( iree_vm_context_t* context, iree_vm_stack_t* stack, iree_vm_function_t function, const iree_vm_invocation_policy_t* policy, - iree_vm_variant_list_t* inputs, iree_vm_variant_list_t* outputs) { + iree_vm_list_t* inputs, iree_vm_list_t* outputs) { // TODO(#2075): disabled because check_test is invoking native methods. // These checks should be nice and simple as we don't support variadic // args/results in bytecode. - iree_host_size_t input_count = inputs ? iree_vm_variant_list_size(inputs) : 0; - iree_host_size_t output_count = - outputs ? iree_vm_variant_list_capacity(outputs) : 0; + iree_host_size_t input_count = inputs ? iree_vm_list_size(inputs) : 0; + iree_host_size_t output_count = outputs ? iree_vm_list_capacity(outputs) : 0; // iree_vm_function_signature_t signature = // iree_vm_function_signature(&function); // if (input_count != signature.argument_count) { @@ -112,8 +111,8 @@ // Marshal inputs into the external stack frame registers. if (inputs) { - iree_vm_stack_frame_marshal_inputs(inputs, external_frame->registers, - argument_registers); + IREE_RETURN_IF_ERROR(iree_vm_stack_frame_marshal_inputs( + inputs, external_frame->registers, argument_registers)); } // Perform execution. Note that for synchronous execution we expect this to @@ -139,8 +138,8 @@ IREE_API_EXPORT iree_status_t IREE_API_CALL iree_vm_invoke( iree_vm_context_t* context, iree_vm_function_t function, - const iree_vm_invocation_policy_t* policy, iree_vm_variant_list_t* inputs, - iree_vm_variant_list_t* outputs, iree_allocator_t allocator) { + const iree_vm_invocation_policy_t* policy, iree_vm_list_t* inputs, + iree_vm_list_t* outputs, iree_allocator_t allocator) { IREE_TRACE_ZONE_BEGIN(z0); // Allocate a VM stack on the host stack and initialize it.
diff --git a/iree/vm/invocation.h b/iree/vm/invocation.h index fc80156..ba62172 100644 --- a/iree/vm/invocation.h +++ b/iree/vm/invocation.h
@@ -19,8 +19,8 @@ #include "iree/base/api.h" #include "iree/vm/context.h" +#include "iree/vm/list.h" #include "iree/vm/module.h" -#include "iree/vm/variant_list.h" #ifdef __cplusplus extern "C" { @@ -46,15 +46,14 @@ // caller. IREE_API_EXPORT iree_status_t IREE_API_CALL iree_vm_invoke( iree_vm_context_t* context, iree_vm_function_t function, - const iree_vm_invocation_policy_t* policy, iree_vm_variant_list_t* inputs, - iree_vm_variant_list_t* outputs, iree_allocator_t allocator); + const iree_vm_invocation_policy_t* policy, iree_vm_list_t* inputs, + iree_vm_list_t* outputs, iree_allocator_t allocator); // TODO(benvanik): document and implement. IREE_API_EXPORT iree_status_t IREE_API_CALL iree_vm_invocation_create( iree_vm_context_t* context, iree_vm_function_t function, - const iree_vm_invocation_policy_t* policy, - const iree_vm_variant_list_t* inputs, iree_allocator_t allocator, - iree_vm_invocation_t** out_invocation); + const iree_vm_invocation_policy_t* policy, const iree_vm_list_t* inputs, + iree_allocator_t allocator, iree_vm_invocation_t** out_invocation); // Retains the given |invocation| for the caller. IREE_API_EXPORT iree_status_t IREE_API_CALL @@ -80,7 +79,7 @@ // released. // // Returns NULL if the invocation did not complete successfully. -IREE_API_EXPORT const iree_vm_variant_list_t* IREE_API_CALL +IREE_API_EXPORT const iree_vm_list_t* IREE_API_CALL iree_vm_invocation_output(iree_vm_invocation_t* invocation); // Blocks the caller until the invocation completes (successfully or otherwise).
diff --git a/iree/vm/list.c b/iree/vm/list.c index 237f506..ac1785e 100644 --- a/iree/vm/list.c +++ b/iree/vm/list.c
@@ -66,38 +66,6 @@ IREE_VM_DEFINE_TYPE_ADAPTERS(iree_vm_list, iree_vm_list_t); -IREE_API_EXPORT iree_status_t IREE_API_CALL iree_vm_list_create( - const iree_vm_type_def_t* element_type, iree_host_size_t initial_capacity, - iree_allocator_t allocator, iree_vm_list_t** out_list) { - iree_vm_list_t* list = NULL; - IREE_RETURN_IF_ERROR( - iree_allocator_malloc(allocator, sizeof(iree_vm_list_t), (void**)&list)); - memset(list, 0, sizeof(*list)); - iree_atomic_store(&list->ref_object.counter, 1); - list->allocator = allocator; - list->element_type = *element_type; - - if (iree_vm_type_def_is_value(element_type)) { - list->storage_mode = IREE_VM_LIST_STORAGE_MODE_VALUE; - list->element_size = kValueTypeSizes[element_type->value_type]; - } else if (iree_vm_type_def_is_ref(element_type)) { - list->storage_mode = IREE_VM_LIST_STORAGE_MODE_REF; - list->element_size = sizeof(iree_vm_ref_t); - } else { - list->storage_mode = IREE_VM_LIST_STORAGE_MODE_VARIANT; - list->element_size = sizeof(iree_vm_variant2_t); - } - - iree_status_t status = iree_vm_list_reserve(list, initial_capacity); - if (!iree_status_is_ok(status)) { - iree_allocator_free(allocator, list); - return status; - } - - *out_list = list; - return IREE_STATUS_OK; -} - static void iree_vm_list_reset_range(iree_vm_list_t* list, iree_host_size_t offset, iree_host_size_t length) { @@ -113,7 +81,7 @@ break; } case IREE_VM_LIST_STORAGE_MODE_VARIANT: { - iree_vm_variant2_t* variant_storage = (iree_vm_variant2_t*)list->storage; + iree_vm_variant_t* variant_storage = (iree_vm_variant_t*)list->storage; for (iree_host_size_t i = offset; i < length; ++i) { if (iree_vm_type_def_is_ref(&variant_storage[i].type)) { iree_vm_ref_release(&variant_storage[i].ref); @@ -124,6 +92,102 @@ } } +IREE_API_EXPORT iree_host_size_t iree_vm_list_storage_size( + const iree_vm_type_def_t* element_type, iree_host_size_t capacity) { + iree_host_size_t element_size = sizeof(iree_vm_variant_t); + if (element_type) { + if (iree_vm_type_def_is_value(element_type)) { + element_size = kValueTypeSizes[element_type->value_type]; + } else if (iree_vm_type_def_is_ref(element_type)) { + element_size = sizeof(iree_vm_ref_t); + } else { + element_size = sizeof(iree_vm_variant_t); + } + } + return iree_align(sizeof(iree_vm_list_t), 8) + + iree_align(capacity * element_size, 8); +} + +IREE_API_EXPORT iree_status_t IREE_API_CALL iree_vm_list_initialize( + iree_byte_span_t storage, const iree_vm_type_def_t* element_type, + iree_host_size_t capacity, iree_vm_list_t** out_list) { + iree_vm_list_storage_mode_t storage_mode = IREE_VM_LIST_STORAGE_MODE_VARIANT; + iree_host_size_t element_size = sizeof(iree_vm_variant_t); + if (element_type) { + if (iree_vm_type_def_is_value(element_type)) { + storage_mode = IREE_VM_LIST_STORAGE_MODE_VALUE; + element_size = kValueTypeSizes[element_type->value_type]; + } else if (iree_vm_type_def_is_ref(element_type)) { + storage_mode = IREE_VM_LIST_STORAGE_MODE_REF; + element_size = sizeof(iree_vm_ref_t); + } else { + storage_mode = IREE_VM_LIST_STORAGE_MODE_VARIANT; + element_size = sizeof(iree_vm_variant_t); + } + } + + iree_host_size_t storage_offset = iree_align(sizeof(iree_vm_list_t), 8); + iree_host_size_t required_storage_size = + storage_offset + iree_align(capacity * element_size, 8); + if (storage.data_length < required_storage_size) { + return IREE_STATUS_OUT_OF_RANGE; + } + memset(storage.data, 0, required_storage_size); + + iree_vm_list_t* list = (iree_vm_list_t*)storage.data; + iree_atomic_store(&list->ref_object.counter, 1); + if (element_type) { + list->element_type = *element_type; + } + list->element_size = element_size; + list->storage_mode = storage_mode; + list->capacity = capacity; + list->storage = storage.data + storage_offset; + + *out_list = list; + return IREE_STATUS_OK; +} + +IREE_API_EXPORT void IREE_API_CALL +iree_vm_list_deinitialize(iree_vm_list_t* list) { + iree_vm_list_reset_range(list, 0, list->count); + list->count = 0; +} + +IREE_API_EXPORT iree_status_t IREE_API_CALL iree_vm_list_create( + const iree_vm_type_def_t* element_type, iree_host_size_t initial_capacity, + iree_allocator_t allocator, iree_vm_list_t** out_list) { + iree_vm_list_t* list = NULL; + IREE_RETURN_IF_ERROR( + iree_allocator_malloc(allocator, sizeof(iree_vm_list_t), (void**)&list)); + memset(list, 0, sizeof(*list)); + iree_atomic_store(&list->ref_object.counter, 1); + list->allocator = allocator; + if (element_type) { + list->element_type = *element_type; + } + + if (iree_vm_type_def_is_value(&list->element_type)) { + list->storage_mode = IREE_VM_LIST_STORAGE_MODE_VALUE; + list->element_size = kValueTypeSizes[element_type->value_type]; + } else if (iree_vm_type_def_is_ref(&list->element_type)) { + list->storage_mode = IREE_VM_LIST_STORAGE_MODE_REF; + list->element_size = sizeof(iree_vm_ref_t); + } else { + list->storage_mode = IREE_VM_LIST_STORAGE_MODE_VARIANT; + list->element_size = sizeof(iree_vm_variant_t); + } + + iree_status_t status = iree_vm_list_reserve(list, initial_capacity); + if (!iree_status_is_ok(status)) { + iree_allocator_free(allocator, list); + return status; + } + + *out_list = list; + return IREE_STATUS_OK; +} + static void iree_vm_list_destroy(void* ptr) { iree_vm_list_t* list = (iree_vm_list_t*)ptr; iree_vm_list_reset_range(list, 0, list->count); @@ -284,12 +348,13 @@ break; } case IREE_VM_LIST_STORAGE_MODE_VARIANT: { - iree_vm_variant2_t* variant = (iree_vm_variant2_t*)element_ptr; + iree_vm_variant_t* variant = (iree_vm_variant_t*)element_ptr; if (!iree_vm_type_def_is_value(&variant->type)) { return IREE_STATUS_FAILED_PRECONDITION; } out_value->type = variant->type.value_type; - out_value->i64 = variant->i64; + memcpy(out_value->value_storage, variant->value_storage, + sizeof(out_value->value_storage)); break; } default: @@ -326,12 +391,13 @@ break; } case IREE_VM_LIST_STORAGE_MODE_VARIANT: { - iree_vm_variant2_t* variant = (iree_vm_variant2_t*)element_ptr; + iree_vm_variant_t* variant = (iree_vm_variant_t*)element_ptr; if (!iree_vm_type_def_is_value(&variant->type)) { return IREE_STATUS_FAILED_PRECONDITION; } value.type = variant->type.value_type; - value.i64 = variant->i64; + memcpy(value.value_storage, variant->value_storage, + sizeof(value.value_storage)); break; } default: @@ -380,10 +446,58 @@ break; } case IREE_VM_LIST_STORAGE_MODE_VARIANT: { - iree_vm_variant2_t* variant = (iree_vm_variant2_t*)element_ptr; + iree_vm_variant_t* variant = (iree_vm_variant_t*)element_ptr; + if (variant->type.ref_type) { + iree_vm_ref_release(&variant->ref); + } variant->type.value_type = target_type; variant->type.ref_type = IREE_VM_REF_TYPE_NULL; - variant->i64 = converted_value.i64; + memcpy(variant->value_storage, converted_value.value_storage, + sizeof(variant->value_storage)); + break; + } + default: + return IREE_STATUS_FAILED_PRECONDITION; + } + return IREE_STATUS_OK; +} + +IREE_API_EXPORT iree_status_t IREE_API_CALL +iree_vm_list_push_value(iree_vm_list_t* list, const iree_vm_value_t* value) { + iree_host_size_t i = iree_vm_list_size(list); + IREE_RETURN_IF_ERROR(iree_vm_list_resize(list, i + 1)); + return iree_vm_list_set_value(list, i, value); +} + +IREE_API_EXPORT void* iree_vm_list_get_ref_deref( + const iree_vm_list_t* list, iree_host_size_t i, + const iree_vm_ref_type_descriptor_t* type_descriptor) { + iree_vm_ref_t value = {0}; + if (!iree_status_is_ok(iree_vm_list_get_ref_assign(list, i, &value))) { + return NULL; + } else if (!iree_status_is_ok( + iree_vm_ref_check(&value, type_descriptor->type))) { + return NULL; + } + return value.ptr; +} + +IREE_API_EXPORT iree_status_t IREE_API_CALL iree_vm_list_get_ref_assign( + const iree_vm_list_t* list, iree_host_size_t i, iree_vm_ref_t* out_value) { + if (i >= list->count) return IREE_STATUS_OUT_OF_RANGE; + uintptr_t element_ptr = (uintptr_t)list->storage + i * list->element_size; + switch (list->storage_mode) { + case IREE_VM_LIST_STORAGE_MODE_REF: { + iree_vm_ref_t* element_ref = (iree_vm_ref_t*)element_ptr; + iree_vm_ref_assign(element_ref, out_value); + break; + } + case IREE_VM_LIST_STORAGE_MODE_VARIANT: { + iree_vm_variant_t* variant = (iree_vm_variant_t*)element_ptr; + if (!iree_vm_type_def_is_ref(&variant->type)) { + return IREE_STATUS_FAILED_PRECONDITION; + } + iree_vm_ref_assign(&variant->ref, out_value); break; } default: @@ -394,25 +508,8 @@ IREE_API_EXPORT iree_status_t IREE_API_CALL iree_vm_list_get_ref_retain( const iree_vm_list_t* list, iree_host_size_t i, iree_vm_ref_t* out_value) { - if (i >= list->count) return IREE_STATUS_OUT_OF_RANGE; - uintptr_t element_ptr = (uintptr_t)list->storage + i * list->element_size; - switch (list->storage_mode) { - case IREE_VM_LIST_STORAGE_MODE_REF: { - iree_vm_ref_t* element_ref = (iree_vm_ref_t*)element_ptr; - iree_vm_ref_retain(element_ref, out_value); - break; - } - case IREE_VM_LIST_STORAGE_MODE_VARIANT: { - iree_vm_variant2_t* variant = (iree_vm_variant2_t*)element_ptr; - if (!iree_vm_type_def_is_ref(&variant->type)) { - return IREE_STATUS_FAILED_PRECONDITION; - } - iree_vm_ref_retain(&variant->ref, out_value); - break; - } - default: - return IREE_STATUS_FAILED_PRECONDITION; - } + IREE_RETURN_IF_ERROR(iree_vm_list_get_ref_assign(list, i, out_value)); + iree_vm_ref_retain(out_value, out_value); return IREE_STATUS_OK; } @@ -430,7 +527,10 @@ break; } case IREE_VM_LIST_STORAGE_MODE_VARIANT: { - iree_vm_variant2_t* variant = (iree_vm_variant2_t*)element_ptr; + iree_vm_variant_t* variant = (iree_vm_variant_t*)element_ptr; + if (variant->type.value_type) { + memset(&variant->ref, 0, sizeof(variant->ref)); + } variant->type.value_type = IREE_VM_VALUE_TYPE_NONE; variant->type.ref_type = value->type; iree_vm_ref_retain_or_move(is_move, value, &variant->ref); @@ -448,22 +548,72 @@ (iree_vm_ref_t*)value); } +IREE_API_EXPORT iree_status_t IREE_API_CALL +iree_vm_list_push_ref_retain(iree_vm_list_t* list, const iree_vm_ref_t* value) { + iree_host_size_t i = iree_vm_list_size(list); + IREE_RETURN_IF_ERROR(iree_vm_list_resize(list, i + 1)); + return iree_vm_list_set_ref_retain(list, i, value); +} + IREE_API_EXPORT iree_status_t IREE_API_CALL iree_vm_list_set_ref_move( iree_vm_list_t* list, iree_host_size_t i, iree_vm_ref_t* value) { return iree_vm_list_set_ref(list, i, /*is_move=*/true, value); } IREE_API_EXPORT iree_status_t IREE_API_CALL +iree_vm_list_push_ref_move(iree_vm_list_t* list, iree_vm_ref_t* value) { + iree_host_size_t i = iree_vm_list_size(list); + IREE_RETURN_IF_ERROR(iree_vm_list_resize(list, i + 1)); + return iree_vm_list_set_ref_move(list, i, value); +} + +IREE_API_EXPORT iree_status_t IREE_API_CALL iree_vm_list_get_variant(const iree_vm_list_t* list, iree_host_size_t i, - iree_vm_variant2_t* out_value) { - return IREE_STATUS_UNIMPLEMENTED; + iree_vm_variant_t* out_value) { + if (i >= list->count) return IREE_STATUS_OUT_OF_RANGE; + uintptr_t element_ptr = (uintptr_t)list->storage + i * list->element_size; + switch (list->storage_mode) { + case IREE_VM_LIST_STORAGE_MODE_VALUE: { + out_value->type = list->element_type; + memcpy(out_value->value_storage, (void*)element_ptr, list->element_size); + break; + } + case IREE_VM_LIST_STORAGE_MODE_REF: { + iree_vm_ref_t* element_ref = (iree_vm_ref_t*)element_ptr; + out_value->type.ref_type = element_ref->type; + out_value->type.value_type = IREE_VM_VALUE_TYPE_NONE; + iree_vm_ref_retain(element_ref, &out_value->ref); + break; + } + case IREE_VM_LIST_STORAGE_MODE_VARIANT: { + iree_vm_variant_t* variant = (iree_vm_variant_t*)element_ptr; + out_value->type = variant->type; + if (iree_vm_type_def_is_ref(&variant->type)) { + iree_vm_ref_assign(&variant->ref, &out_value->ref); + } else { + memcpy(out_value->value_storage, variant->value_storage, + sizeof(variant->value_storage)); + } + break; + } + default: + return IREE_STATUS_FAILED_PRECONDITION; + } + return IREE_STATUS_OK; } IREE_API_EXPORT iree_status_t IREE_API_CALL iree_vm_list_set_variant( - iree_vm_list_t* list, iree_host_size_t i, const iree_vm_variant2_t* value) { + iree_vm_list_t* list, iree_host_size_t i, const iree_vm_variant_t* value) { return IREE_STATUS_UNIMPLEMENTED; } +IREE_API_EXPORT iree_status_t IREE_API_CALL iree_vm_list_push_variant( + iree_vm_list_t* list, const iree_vm_variant_t* value) { + iree_host_size_t i = iree_vm_list_size(list); + IREE_RETURN_IF_ERROR(iree_vm_list_resize(list, i + 1)); + return iree_vm_list_set_variant(list, i, value); +} + iree_status_t iree_vm_list_register_types() { iree_vm_list_descriptor.destroy = iree_vm_list_destroy; iree_vm_list_descriptor.offsetof_counter =
diff --git a/iree/vm/list.h b/iree/vm/list.h index d0700db..00fd046 100644 --- a/iree/vm/list.h +++ b/iree/vm/list.h
@@ -33,15 +33,46 @@ // ABI. It is not designed for efficiency: if you are performing large amounts // of work on the list type you should instead be representing that using the // HAL types so that you can get acceleration. +// +// This type the same performance characteristics as std::vector; pushes may +// grow the capacity of the list and to ensure minimal wastage it is always +// better to reserve the exact desired element count first. typedef struct iree_vm_list iree_vm_list_t; #ifndef IREE_API_NO_PROTOTYPES +// Returns the size in bytes required to store a list with the given element +// type and capacity. This storage size can be used to stack allocate or reserve +// memory that is then used by iree_vm_list_initialize to avoid dynamic +// allocations. +IREE_API_EXPORT iree_host_size_t iree_vm_list_storage_size( + const iree_vm_type_def_t* element_type, iree_host_size_t capacity); + +// Initializes a statically-allocated list in the |storage| memory. +// The storage capacity must be large enough to hold the list internals and +// its contents which may vary across compilers/platforms/etc; use +// iree_vm_list_storage_size to query the required capacity. +// +// Statically-allocated lists have their lifetime controlled by the caller and +// must be deinitialized with iree_vm_list_deinitialize only when there are no +// more users of the list. +IREE_API_EXPORT iree_status_t IREE_API_CALL iree_vm_list_initialize( + iree_byte_span_t storage, const iree_vm_type_def_t* element_type, + iree_host_size_t capacity, iree_vm_list_t** out_list); + +// Deinitializes a statically-allocated |list| previously initialized with +// iree_vm_list_initialize. +IREE_API_EXPORT void IREE_API_CALL +iree_vm_list_deinitialize(iree_vm_list_t* list); + // Creates a growable list containing the given |element_type|, which may either // be a primitive iree_vm_value_type_t value (like i32) or a ref type. When // storing ref types the list may either store a specific iree_vm_ref_type_t // and ensure that all elements set match the type or IREE_VM_REF_TYPE_ANY to // indicate that any ref type is allowed. +// +// |element_type| can be set to iree_vm_type_def_make_variant_type (or null) to +// indicate that the list stores variants (each element can differ in type). IREE_API_EXPORT iree_status_t IREE_API_CALL iree_vm_list_create( const iree_vm_type_def_t* element_type, iree_host_size_t initial_capacity, iree_allocator_t allocator, iree_vm_list_t** out_list); @@ -95,6 +126,25 @@ IREE_API_EXPORT iree_status_t IREE_API_CALL iree_vm_list_set_value( iree_vm_list_t* list, iree_host_size_t i, const iree_vm_value_t* value); +// Pushes the value of the element to the end of the list. +// If the specified |value| type differs from the list storage type the value +// will be converted using the value type semantics (such as sign/zero extend, +// etc). +IREE_API_EXPORT iree_status_t IREE_API_CALL +iree_vm_list_push_value(iree_vm_list_t* list, const iree_vm_value_t* value); + +// Returns a dereferenced pointer to the given type if the element at the given +// index matches the type. Returns NULL on error. +IREE_API_EXPORT void* iree_vm_list_get_ref_deref( + const iree_vm_list_t* list, iree_host_size_t i, + const iree_vm_ref_type_descriptor_t* type_descriptor); + +// Returns the ref value of the element at the given index. +// The ref will not be retained and must be retained by the caller to extend +// its lifetime. +IREE_API_EXPORT iree_status_t IREE_API_CALL iree_vm_list_get_ref_assign( + const iree_vm_list_t* list, iree_host_size_t i, iree_vm_ref_t* out_value); + // Returns the ref value of the element at the given index. // The ref will be retained and must be released by the caller. IREE_API_EXPORT iree_status_t IREE_API_CALL iree_vm_list_get_ref_retain( @@ -105,23 +155,41 @@ IREE_API_EXPORT iree_status_t IREE_API_CALL iree_vm_list_set_ref_retain( iree_vm_list_t* list, iree_host_size_t i, const iree_vm_ref_t* value); +// Pushes the ref value of the element to the end of the list, retaining a +// reference in the list until the element is cleared or the list is disposed. +IREE_API_EXPORT iree_status_t IREE_API_CALL +iree_vm_list_push_ref_retain(iree_vm_list_t* list, const iree_vm_ref_t* value); + // Sets the ref value of the element at the given index, moving ownership of the // |value| reference to the list. IREE_API_EXPORT iree_status_t IREE_API_CALL iree_vm_list_set_ref_move( iree_vm_list_t* list, iree_host_size_t i, iree_vm_ref_t* value); +// Pushes the ref value of the element to the end of the list, moving ownership +// of the |value| reference to the list. +IREE_API_EXPORT iree_status_t IREE_API_CALL +iree_vm_list_push_ref_move(iree_vm_list_t* list, iree_vm_ref_t* value); + // Returns the value of the element at the given index. If the element contains -// a ref then it will be retained and must be released by the caller. +// a ref it will *not* be retained and the caller must retain it to extend its +// lifetime. IREE_API_EXPORT iree_status_t IREE_API_CALL iree_vm_list_get_variant(const iree_vm_list_t* list, iree_host_size_t i, - iree_vm_variant2_t* out_value); + iree_vm_variant_t* out_value); // Sets the value of the element at the given index. If the specified |value| // type differs from the list storage type the value will be converted using the // value type semantics (such as sign/zero extend, etc). If the variant is a ref // then it will be retained. IREE_API_EXPORT iree_status_t IREE_API_CALL iree_vm_list_set_variant( - iree_vm_list_t* list, iree_host_size_t i, const iree_vm_variant2_t* value); + iree_vm_list_t* list, iree_host_size_t i, const iree_vm_variant_t* value); + +// Pushes the value of the element to the end of the list. If the specified +// |value| type differs from the list storage type the value will be converted +// using the value type semantics (such as sign/zero extend, etc). If the +// variant is a ref then it will be retained. +IREE_API_EXPORT iree_status_t IREE_API_CALL +iree_vm_list_push_variant(iree_vm_list_t* list, const iree_vm_variant_t* value); #endif // IREE_API_NO_PROTOTYPES
diff --git a/iree/vm/module.h b/iree/vm/module.h index e278992..3081e10 100644 --- a/iree/vm/module.h +++ b/iree/vm/module.h
@@ -210,7 +210,7 @@ // attributes. // Returns IREE_STATUS_NOT_FOUND if index >= the number of attributes for // the function. - // See: docs/function_abi.md + // See: docs/design_docs/function_abi.md iree_status_t(IREE_API_PTR* get_function_reflection_attr)( void* self, iree_vm_function_linkage_t linkage, int32_t ordinal, int32_t index, iree_string_view_t* key, iree_string_view_t* value); @@ -277,7 +277,7 @@ // Returns the empty string if the reflection data in general or the specific // key is not found. // -// See: docs/function_abi.md for documentation on the ABI. +// See: docs/design_docs/function_abi.md for documentation on the ABI. IREE_API_EXPORT iree_string_view_t IREE_API_CALL iree_vm_function_reflection_attr(const iree_vm_function_t* function, iree_string_view_t key); @@ -289,7 +289,7 @@ // attributes. // Returns IREE_STATUS_NOT_FOUND if index >= the number of attributes for // the function. -// See: docs/function_abi.md +// See: docs/design_docs/function_abi.md IREE_API_EXPORT iree_status_t IREE_API_CALL iree_vm_get_function_reflection_attr(iree_vm_function_t function, int32_t index, iree_string_view_t* key,
diff --git a/iree/vm/ref_cc.h b/iree/vm/ref_cc.h index 8120f24..13006d1 100644 --- a/iree/vm/ref_cc.h +++ b/iree/vm/ref_cc.h
@@ -107,39 +107,55 @@ return ref_type_descriptor<T>::get()->type; } - ABSL_ATTRIBUTE_ALWAYS_INLINE ref() noexcept = default; - ABSL_ATTRIBUTE_ALWAYS_INLINE ref(std::nullptr_t) noexcept {} // NOLINT - ABSL_ATTRIBUTE_ALWAYS_INLINE ref(T* p) noexcept : px_(p) {} - ABSL_ATTRIBUTE_ALWAYS_INLINE ~ref() noexcept { - if (px_) ref_type_release<T>(px_); - } + ABSL_ATTRIBUTE_ALWAYS_INLINE ref() noexcept + : ref_({ + 0, + ref_type_descriptor<T>::get()->offsetof_counter, + ref_type_descriptor<T>::get()->type, + }) {} + ABSL_ATTRIBUTE_ALWAYS_INLINE ref(std::nullptr_t) noexcept + : ref_({ + 0, + ref_type_descriptor<T>::get()->offsetof_counter, + ref_type_descriptor<T>::get()->type, + }) {} + ABSL_ATTRIBUTE_ALWAYS_INLINE ref(T* p) noexcept + : ref_({ + p, + ref_type_descriptor<T>::get()->offsetof_counter, + ref_type_descriptor<T>::get()->type, + }) {} + ABSL_ATTRIBUTE_ALWAYS_INLINE ~ref() noexcept { ref_type_release<T>(get()); } // Don't use implicit ref copying; use retain_ref instead to make things more // readable. We can't delete the ctor (or, I couldn't find a way not to) // because the templated parameter packing magic needs it. - ref(const ref& rhs) noexcept : px_(rhs.get()) { - if (px_) ref_type_retain<T>(px_); - } + ref(const ref& rhs) noexcept : ref_(rhs.ref_) { ref_type_retain<T>(get()); } ref& operator=(const ref&) noexcept = delete; // Move support to transfer ownership from one ref to another. - ref(ref&& rhs) noexcept : px_(rhs.release()) {} + ref(ref&& rhs) noexcept : ref_(rhs.ref_) { rhs.release(); } ref& operator=(ref&& rhs) noexcept { - if (px_ != rhs.px_) { - if (px_) ref_type_release<T>(px_); - px_ = rhs.release(); + if (get() != rhs.get()) { + ref_type_release<T>(get()); + ref_ = rhs.ref_; + rhs.release(); } return *this; } // Move support from another compatible type. template <typename U> - ref(ref<U>&& rhs) noexcept : px_(rhs.release()) {} // NOLINT + ref(ref<U>&& rhs) noexcept { + ref_.ptr = static_cast<T*>(rhs.release()); + ref_.offsetof_counter = rhs.ref_.offsetof_counter; + ref_.type = rhs.ref_.type; + } template <typename U> ref& operator=(ref<U>&& rhs) noexcept { - if (px_ != rhs.get()) { - if (px_) ref_type_release<T>(px_); - px_ = rhs.release(); + if (get() != rhs.get()) { + ref_type_release<T>(get()); + ref_.ptr = static_cast<T*>(rhs.release()); } return *this; } @@ -147,10 +163,8 @@ // Resets the object to nullptr and decrements the reference count, possibly // deleting it. void reset() noexcept { - if (px_) { - ref_type_release<T>(px_); - px_ = nullptr; - } + ref_type_release<T>(get()); + ref_.ptr = nullptr; } // Releases a pointer. @@ -159,8 +173,8 @@ // Returns nullptr if the ref holds no value. // To re-wrap in a ref use either ref<T>(value) or assign(). ABSL_ATTRIBUTE_ALWAYS_INLINE T* release() noexcept { - T* p = px_; - px_ = nullptr; + T* p = get(); + ref_.ptr = nullptr; return p; } @@ -169,33 +183,42 @@ // not be incremented. ABSL_ATTRIBUTE_ALWAYS_INLINE void assign(T* value) noexcept { reset(); - px_ = value; + ref_.ptr = value; } // Gets the pointer referenced by this instance. // operator* and operator-> will assert() if there is no current object. - constexpr T* get() const noexcept { return px_; } - constexpr T& operator*() const noexcept { return *px_; } - constexpr T* operator->() const noexcept { return px_; } + constexpr T* get() const noexcept { return reinterpret_cast<T*>(ref_.ptr); } + constexpr T& operator*() const noexcept { return *get(); } + constexpr T* operator->() const noexcept { return get(); } // Returns a pointer to the inner pointer storage. // This allows passing a pointer to the ref as an output argument to C-style // creation functions. - constexpr T** operator&() noexcept { return &px_; } // NOLINT + constexpr T** operator&() noexcept { + return reinterpret_cast<T**>(&ref_.ptr); + } // Support boolean expression evaluation ala unique_ptr/shared_ptr: // https://en.cppreference.com/w/cpp/memory/shared_ptr/operator_bool constexpr operator unspecified_bool_type() const noexcept { - return px_ ? &this_type::px_ : nullptr; + return get() ? reinterpret_cast<unspecified_bool_type>(&this_type::ref_.ptr) + : nullptr; } // Supports unary expression evaluation. - constexpr bool operator!() const noexcept { return !px_; } + constexpr bool operator!() const noexcept { return !get(); } // Swap support. - void swap(ref& rhs) { std::swap(px_, rhs.px_); } + void swap(ref& rhs) { std::swap(ref_.ptr, rhs.ref_.ptr); } + + // Allows directly passing the ref to a C-API function for creation. + // Example: + // iree::vm::ref<my_type_t> value; + // my_type_create(..., &value); + constexpr operator iree_vm_ref_t*() const noexcept { return &ref_; } private: - T* px_ = nullptr; + mutable iree_vm_ref_t ref_; }; // Adds a reference to the given ref and returns the same ref. @@ -206,7 +229,7 @@ // retain_ref(b); // ref count + 1 template <typename T> inline ref<T> retain_ref(const ref<T>& value) { - if (value) ref_type_retain<T>(value.get()); + ref_type_retain<T>(value.get()); return ref<T>(value.get()); } @@ -217,7 +240,7 @@ // ref<MyType> p = retain_ref(raw_ptr); // ref count + 1 template <typename T> inline ref<T> retain_ref(T* value) { - if (value) ref_type_retain<T>(value); + ref_type_retain<T>(value); return ref<T>(value); }
diff --git a/iree/vm/stack.c b/iree/vm/stack.c index dc4e405..d65659c 100644 --- a/iree/vm/stack.c +++ b/iree/vm/stack.c
@@ -499,8 +499,8 @@ const iree_vm_register_list_t* src_reg_list, const iree_vm_registers_t dst_regs, const iree_vm_register_list_t* dst_reg_list) { - VMCHECK(src_reg_list->size == dst_reg_list->size); - if (src_reg_list->size != dst_reg_list->size) return; + VMCHECK(src_reg_list->size <= dst_reg_list->size); + if (src_reg_list->size > dst_reg_list->size) return; for (int i = 0; i < src_reg_list->size; ++i) { // TODO(benvanik): change encoding to avoid this branching. // Could write two arrays: one for prims and one for refs. @@ -637,6 +637,7 @@ VMCHECK(external_registers->size >= callee_registers->size); uint16_t i32_reg_ordinal = 0; uint16_t ref_reg_ordinal = 0; + ((iree_vm_register_list_t*)external_registers)->size = callee_registers->size; uint16_t* dst_reg_list = (uint16_t*)external_registers->registers; for (int i = 0; i < callee_registers->size; ++i) { uint16_t src_reg = callee_registers->registers[i];
diff --git a/iree/vm/stack.h b/iree/vm/stack.h index ffbc1e0..7bd3535 100644 --- a/iree/vm/stack.h +++ b/iree/vm/stack.h
@@ -22,7 +22,6 @@ #include "iree/base/api.h" #include "iree/vm/module.h" #include "iree/vm/ref.h" -#include "iree/vm/variant_list.h" #ifdef __cplusplus extern "C" {
diff --git a/iree/vm/test/BUILD b/iree/vm/test/BUILD index 5018865..c33c07d 100644 --- a/iree/vm/test/BUILD +++ b/iree/vm/test/BUILD
@@ -25,6 +25,7 @@ name = "all_bytecode_modules_cc", srcs = [ ":arithmetic_ops.module", + ":comparison_ops.module", ":control_flow_ops.module", ":list_ops.module", ], @@ -41,6 +42,12 @@ ) iree_bytecode_module( + name = "comparison_ops", + src = "comparison_ops.mlir", + flags = ["-iree-vm-ir-to-bytecode-module"], +) + +iree_bytecode_module( name = "control_flow_ops", src = "control_flow_ops.mlir", flags = ["-iree-vm-ir-to-bytecode-module"],
diff --git a/iree/vm/test/CMakeLists.txt b/iree/vm/test/CMakeLists.txt index cb6951f..747595b 100644 --- a/iree/vm/test/CMakeLists.txt +++ b/iree/vm/test/CMakeLists.txt
@@ -19,6 +19,7 @@ all_bytecode_modules_cc GENERATED_SRCS "arithmetic_ops.module" + "comparison_ops.module" "control_flow_ops.module" "list_ops.module" CC_FILE_OUTPUT @@ -43,6 +44,16 @@ iree_bytecode_module( NAME + comparison_ops + SRC + "comparison_ops.mlir" + FLAGS + "-iree-vm-ir-to-bytecode-module" + PUBLIC +) + +iree_bytecode_module( + NAME control_flow_ops SRC "control_flow_ops.mlir"
diff --git a/iree/vm/test/comparison_ops.mlir b/iree/vm/test/comparison_ops.mlir new file mode 100644 index 0000000..7161cb6 --- /dev/null +++ b/iree/vm/test/comparison_ops.mlir
@@ -0,0 +1,172 @@ +vm.module @comparison_ops { + + //===--------------------------------------------------------------------===// + // vm.cmp.lt.i32.s + //===--------------------------------------------------------------------===// + + vm.export @test_cmp_lt_s_0 + vm.func @test_cmp_lt_s_0() { + %lhs = vm.const.i32 2 : i32 + %lhs_dno = iree.do_not_optimize(%lhs) : i32 + %rhs = vm.const.i32 -2 : i32 + %rhs_dno = iree.do_not_optimize(%rhs) : i32 + %actual = vm.cmp.lt.i32.s %lhs_dno, %rhs_dno : i32 + %expected = vm.const.i32 0 : i32 + vm.check.eq %actual, %expected, "2 < -2" : i32 + vm.return + } + + vm.export @test_cmp_lt_s_1 + vm.func @test_cmp_lt_s_1() { + %lhs = vm.const.i32 -2 : i32 + %lhs_dno = iree.do_not_optimize(%lhs) : i32 + %rhs = vm.const.i32 2 : i32 + %rhs_dno = iree.do_not_optimize(%rhs) : i32 + %actual = vm.cmp.lt.i32.s %lhs_dno, %rhs_dno : i32 + %expected = vm.const.i32 1 : i32 + vm.check.eq %actual, %expected, "-2 < 2" : i32 + vm.return + } + + // Expect UINT_MAX to be interpreted as -1 when doing a signed compare. + vm.export @test_cmp_lt_s_2 + vm.func @test_cmp_lt_s_2() { + %lhs = vm.const.i32 4294967295 : i32 + %lhs_dno = iree.do_not_optimize(%lhs) : i32 + %rhs = vm.const.i32 2 : i32 + %rhs_dno = iree.do_not_optimize(%rhs) : i32 + %actual = vm.cmp.lt.i32.s %lhs_dno, %rhs_dno : i32 + %expected = vm.const.i32 1 : i32 + vm.check.eq %actual, %expected, "4294967295 (UINT_MAX) < 2" : i32 + vm.return + } + + //===--------------------------------------------------------------------===// + // vm.cmp.lt.i32.u + //===--------------------------------------------------------------------===// + + vm.export @test_cmp_lt_u_0 + vm.func @test_cmp_lt_u_0() { + %lhs = vm.const.i32 2 : i32 + %lhs_dno = iree.do_not_optimize(%lhs) : i32 + %rhs = vm.const.i32 -2 : i32 + %rhs_dno = iree.do_not_optimize(%rhs) : i32 + %actual = vm.cmp.lt.i32.u %lhs_dno, %rhs_dno : i32 + %expected = vm.const.i32 1 : i32 + vm.check.eq %actual, %expected, "2 < -2 (as unsigned)" : i32 + vm.return + } + + vm.export @test_cmp_lt_u_1 + vm.func @test_cmp_lt_u_1() { + %lhs = vm.const.i32 -2 : i32 + %lhs_dno = iree.do_not_optimize(%lhs) : i32 + %rhs = vm.const.i32 2 : i32 + %rhs_dno = iree.do_not_optimize(%rhs) : i32 + %actual = vm.cmp.lt.i32.u %lhs_dno, %rhs_dno : i32 + %expected = vm.const.i32 0 : i32 + vm.check.eq %actual, %expected, "-2 < 2 (as unsigned)" : i32 + vm.return + } + + vm.export @test_cmp_lt_u_2 + vm.func @test_cmp_lt_u_2() { + %lhs = vm.const.i32 4294967295 : i32 + %lhs_dno = iree.do_not_optimize(%lhs) : i32 + %rhs = vm.const.i32 2 : i32 + %rhs_dno = iree.do_not_optimize(%rhs) : i32 + %actual = vm.cmp.lt.i32.u %lhs_dno, %rhs_dno : i32 + %expected = vm.const.i32 0 : i32 + vm.check.eq %actual, %expected, "4294967295 (UINT_MAX) < 2 (as unsigned)" : i32 + vm.return + } + + //===--------------------------------------------------------------------===// + // vm.cmp.*.i32.* pseudo-ops + //===--------------------------------------------------------------------===// + // NOTE: all of these are turned in to some variants of vm.cmp.lt by the + // compiler and are here as a way to test the runtime behavior of the + // pseudo-op expansions. + + vm.export @test_cmp_lte + vm.func @test_cmp_lte() { + %true = vm.const.i32 1 : i32 + %false = vm.const.i32 0 : i32 + + %cn2 = vm.const.i32 -2 : i32 + %cn2_dno = iree.do_not_optimize(%cn2) : i32 + %c2 = vm.const.i32 2 : i32 + %c2_dno = iree.do_not_optimize(%c2) : i32 + + %cmp_0 = vm.cmp.lte.i32.s %cn2_dno, %c2_dno : i32 + vm.check.eq %cmp_0, %true, "-2 <= 2" : i32 + %cmp_1 = vm.cmp.lte.i32.s %c2_dno, %cn2_dno : i32 + vm.check.eq %cmp_1, %false, "2 <= -2" : i32 + %cmp_2 = vm.cmp.lte.i32.s %c2_dno, %c2_dno : i32 + vm.check.eq %cmp_2, %true, "2 <= 2" : i32 + + %cmp_3 = vm.cmp.lte.i32.u %cn2_dno, %c2_dno : i32 + vm.check.eq %cmp_3, %false, "-2 <= 2 (unsigned)" : i32 + %cmp_4 = vm.cmp.lte.i32.u %c2_dno, %cn2_dno : i32 + vm.check.eq %cmp_4, %true, "2 <= -2 (unsigned)" : i32 + %cmp_5 = vm.cmp.lte.i32.u %c2_dno, %c2_dno : i32 + vm.check.eq %cmp_5, %true, "2 <= 2 (unsigned)" : i32 + + vm.return + } + + vm.export @test_cmp_gt + vm.func @test_cmp_gt() { + %true = vm.const.i32 1 : i32 + %false = vm.const.i32 0 : i32 + + %cn2 = vm.const.i32 -2 : i32 + %cn2_dno = iree.do_not_optimize(%cn2) : i32 + %c2 = vm.const.i32 2 : i32 + %c2_dno = iree.do_not_optimize(%c2) : i32 + + %cmp_0 = vm.cmp.gt.i32.s %cn2_dno, %c2_dno : i32 + vm.check.eq %cmp_0, %false, "-2 > 2" : i32 + %cmp_1 = vm.cmp.gt.i32.s %c2_dno, %cn2_dno : i32 + vm.check.eq %cmp_1, %true, "2 > -2" : i32 + %cmp_2 = vm.cmp.gt.i32.s %c2_dno, %c2_dno : i32 + vm.check.eq %cmp_2, %false, "2 > 2" : i32 + + %cmp_3 = vm.cmp.gt.i32.u %cn2_dno, %c2_dno : i32 + vm.check.eq %cmp_3, %true, "-2 > 2 (unsigned)" : i32 + %cmp_4 = vm.cmp.gt.i32.u %c2_dno, %cn2_dno : i32 + vm.check.eq %cmp_4, %false, "2 > -2 (unsigned)" : i32 + %cmp_5 = vm.cmp.gt.i32.u %c2_dno, %c2_dno : i32 + vm.check.eq %cmp_5, %false, "2 > 2 (unsigned)" : i32 + + vm.return + } + + vm.export @test_cmp_gte + vm.func @test_cmp_gte() { + %true = vm.const.i32 1 : i32 + %false = vm.const.i32 0 : i32 + + %cn2 = vm.const.i32 -2 : i32 + %cn2_dno = iree.do_not_optimize(%cn2) : i32 + %c2 = vm.const.i32 2 : i32 + %c2_dno = iree.do_not_optimize(%c2) : i32 + + %cmp_0 = vm.cmp.gte.i32.s %cn2_dno, %c2_dno : i32 + vm.check.eq %cmp_0, %false, "-2 >= 2" : i32 + %cmp_1 = vm.cmp.gte.i32.s %c2_dno, %cn2_dno : i32 + vm.check.eq %cmp_1, %true, "2 >= -2" : i32 + %cmp_2 = vm.cmp.gte.i32.s %c2_dno, %c2_dno : i32 + vm.check.eq %cmp_2, %true, "2 >= 2" : i32 + + %cmp_3 = vm.cmp.gte.i32.u %cn2_dno, %c2_dno : i32 + vm.check.eq %cmp_3, %true, "-2 >= 2 (unsigned)" : i32 + %cmp_4 = vm.cmp.gte.i32.u %c2_dno, %cn2_dno : i32 + vm.check.eq %cmp_4, %false, "2 >= -2 (unsigned)" : i32 + %cmp_5 = vm.cmp.gte.i32.u %c2_dno, %c2_dno : i32 + vm.check.eq %cmp_5, %true, "2 >= 2 (unsigned)" : i32 + + vm.return + } + +}
diff --git a/iree/vm/type_def.h b/iree/vm/type_def.h index 4874b3a..2055fa2 100644 --- a/iree/vm/type_def.h +++ b/iree/vm/type_def.h
@@ -84,8 +84,16 @@ int32_t i32; int64_t i64; iree_vm_ref_t ref; + + uint8_t value_storage[IREE_VM_VALUE_STORAGE_SIZE]; // max size of all value + // types }; -} iree_vm_variant2_t; +} iree_vm_variant_t; + +#define iree_vm_variant_empty() \ + { {IREE_VM_VALUE_TYPE_NONE, IREE_VM_REF_TYPE_NULL}, {0}, } +#define iree_vm_variant_is_value(v) iree_vm_type_def_is_value(&v.type) +#define iree_vm_variant_is_ref(v) iree_vm_type_def_is_ref(&v.type) #ifdef __cplusplus } // extern "C"
diff --git a/iree/vm/value.h b/iree/vm/value.h index bb7ecc7..92dc11f 100644 --- a/iree/vm/value.h +++ b/iree/vm/value.h
@@ -38,6 +38,9 @@ IREE_VM_VALUE_TYPE_COUNT = IREE_VM_VALUE_TYPE_MAX + 1, } iree_vm_value_type_t; +// Maximum size, in bytes, of any value type we can represent. +#define IREE_VM_VALUE_STORAGE_SIZE 8 + // A variant value type. typedef struct iree_vm_value { iree_vm_value_type_t type; @@ -46,6 +49,9 @@ int16_t i16; int32_t i32; int64_t i64; + + uint8_t value_storage[IREE_VM_VALUE_STORAGE_SIZE]; // max size of all value + // types }; } iree_vm_value_t;
diff --git a/iree/vm/variant_list.c b/iree/vm/variant_list.c deleted file mode 100644 index 1b6a1e6..0000000 --- a/iree/vm/variant_list.c +++ /dev/null
@@ -1,128 +0,0 @@ -// Copyright 2019 Google LLC -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// https://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. - -#include "iree/vm/variant_list.h" - -struct iree_vm_variant_list { - iree_allocator_t allocator; - iree_host_size_t capacity; - iree_host_size_t count; - iree_vm_variant_t values[]; -}; - -IREE_API_EXPORT iree_status_t IREE_API_CALL iree_vm_variant_list_alloc( - iree_host_size_t capacity, iree_allocator_t allocator, - iree_vm_variant_list_t** out_list) { - if (!out_list) { - return IREE_STATUS_INVALID_ARGUMENT; - } - *out_list = NULL; - - iree_host_size_t alloc_size = iree_vm_variant_list_alloc_size(capacity); - iree_vm_variant_list_t* list = NULL; - IREE_RETURN_IF_ERROR( - iree_allocator_malloc(allocator, alloc_size, (void**)&list)); - iree_vm_variant_list_init(list, capacity); - list->allocator = allocator; - *out_list = list; - return IREE_STATUS_OK; -} - -IREE_API_EXPORT iree_host_size_t IREE_API_CALL -iree_vm_variant_list_alloc_size(iree_host_size_t capacity) { - return sizeof(iree_vm_variant_list_t) + sizeof(iree_vm_variant_t) * capacity; -} - -IREE_API_EXPORT void IREE_API_CALL iree_vm_variant_list_init( - iree_vm_variant_list_t* list, iree_host_size_t capacity) { - memset(&list->allocator, 0, sizeof(list->allocator)); - list->capacity = capacity; - list->count = 0; - memset(list->values, 0, sizeof(list->values[0]) * capacity); -} - -IREE_API_EXPORT void IREE_API_CALL -iree_vm_variant_list_free(iree_vm_variant_list_t* list) { - for (iree_host_size_t i = 0; i < list->count; ++i) { - if (IREE_VM_VARIANT_IS_REF(&list->values[i])) { - iree_vm_ref_release(&list->values[i].ref); - } - } - iree_allocator_free(list->allocator, list); -} - -IREE_API_EXPORT iree_host_size_t IREE_API_CALL -iree_vm_variant_list_capacity(const iree_vm_variant_list_t* list) { - return list->capacity; -} - -IREE_API_EXPORT iree_host_size_t IREE_API_CALL -iree_vm_variant_list_size(const iree_vm_variant_list_t* list) { - return list->count; -} - -IREE_API_EXPORT iree_status_t IREE_API_CALL iree_vm_variant_list_append_value( - iree_vm_variant_list_t* list, iree_vm_value_t value) { - if (list->count + 1 > list->capacity) { - return IREE_STATUS_OUT_OF_RANGE; - } - iree_host_size_t i = list->count++; - list->values[i].value_type = value.type; - list->values[i].ref_type = IREE_VM_REF_TYPE_NULL; - list->values[i].i32 = value.i32; - return IREE_STATUS_OK; -} - -IREE_API_EXPORT iree_status_t IREE_API_CALL -iree_vm_variant_list_append_ref_retain(iree_vm_variant_list_t* list, - iree_vm_ref_t* ref) { - if (list->count + 1 > list->capacity) { - return IREE_STATUS_OUT_OF_RANGE; - } - iree_host_size_t i = list->count++; - list->values[i].value_type = IREE_VM_VALUE_TYPE_NONE; - list->values[i].ref_type = ref->type; - iree_vm_ref_retain(ref, &list->values[i].ref); - return IREE_STATUS_OK; -} - -IREE_API_EXPORT iree_status_t IREE_API_CALL -iree_vm_variant_list_append_ref_move(iree_vm_variant_list_t* list, - iree_vm_ref_t* ref) { - if (list->count + 1 > list->capacity) { - return IREE_STATUS_OUT_OF_RANGE; - } - iree_host_size_t i = list->count++; - list->values[i].value_type = IREE_VM_VALUE_TYPE_NONE; - list->values[i].ref_type = ref->type; - iree_vm_ref_move(ref, &list->values[i].ref); - return IREE_STATUS_OK; -} - -IREE_API_EXPORT iree_status_t IREE_API_CALL -iree_vm_variant_list_append_null_ref(iree_vm_variant_list_t* list) { - if (list->count + 1 > list->capacity) { - return IREE_STATUS_OUT_OF_RANGE; - } - iree_host_size_t i = list->count++; - list->values[i].value_type = IREE_VM_VALUE_TYPE_NONE; - list->values[i].ref_type = IREE_VM_REF_TYPE_NULL; - return IREE_STATUS_OK; -} - -IREE_API_EXPORT iree_vm_variant_t* IREE_API_CALL -iree_vm_variant_list_get(iree_vm_variant_list_t* list, iree_host_size_t i) { - if (i < 0 || i > list->count) return NULL; - return &list->values[i]; -}
diff --git a/iree/vm/variant_list.h b/iree/vm/variant_list.h deleted file mode 100644 index 830efb8..0000000 --- a/iree/vm/variant_list.h +++ /dev/null
@@ -1,111 +0,0 @@ -// Copyright 2019 Google LLC -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// https://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. - -#ifndef IREE_VM_VARIANT_LIST_H_ -#define IREE_VM_VARIANT_LIST_H_ - -#include <stdint.h> - -#include "iree/vm/ref.h" -#include "iree/vm/value.h" - -#ifdef __cplusplus -extern "C" { -#endif // __cplusplus - -// A list able to hold both value and ref types, used for marshaling host -// language calls to VM calls. -// -// Lists may either be stack or heap allocated depending on how long they must -// live. Prefer stack allocations when APIs receiving the lists document -// themselves as copying the list values. -typedef struct iree_vm_variant_list iree_vm_variant_list_t; - -// An element of an iree_vm_variant_list_t. -typedef struct { - iree_vm_value_type_t value_type : 8; - iree_vm_ref_type_t ref_type : 24; - union { - int32_t i32; - iree_vm_ref_t ref; - }; -} iree_vm_variant_t; - -#define IREE_VM_VARIANT_IS_VALUE(v) ((v)->value_type != IREE_VM_VALUE_TYPE_NONE) -#define IREE_VM_VARIANT_IS_REF(v) !IREE_VM_VARIANT_IS_VALUE(v) - -#ifndef IREE_API_NO_PROTOTYPES - -// Allocates a list with the maximum |capacity|. -// The list must be freed with iree_vm_variant_list_free unless ownership is -// transferred to code that will perform the free as documented in its API. -IREE_API_EXPORT iree_status_t IREE_API_CALL iree_vm_variant_list_alloc( - iree_host_size_t capacity, iree_allocator_t allocator, - iree_vm_variant_list_t** out_list); - -// Returns the size, in bytes, required to store a list of the given capacity. -// This can be used to stack-allocate the variant list and then wrap the memory -// for use as a list with iree_vm_variant_list_init. -IREE_API_EXPORT iree_host_size_t IREE_API_CALL -iree_vm_variant_list_alloc_size(iree_host_size_t capacity); - -// Initializes the allocated |list| with the given |capacity|. The list -// allocation must be at least the size returned by -// iree_vm_variant_list_alloc_size for the same |capacity|. -// The list must be freed with iree_vm_variant_list_free unless ownership is -// transferred to code that will perform the free as documented in its API. -IREE_API_EXPORT void IREE_API_CALL iree_vm_variant_list_init( - iree_vm_variant_list_t* list, iree_host_size_t capacity); - -// Frees the list using the allocator it was originally allocated from. -IREE_API_EXPORT void IREE_API_CALL -iree_vm_variant_list_free(iree_vm_variant_list_t* list); - -// Returns the capacity of the list in elements. -IREE_API_EXPORT iree_host_size_t IREE_API_CALL -iree_vm_variant_list_capacity(const iree_vm_variant_list_t* list); - -// Returns the total number of elements added to the list. -IREE_API_EXPORT iree_host_size_t IREE_API_CALL -iree_vm_variant_list_size(const iree_vm_variant_list_t* list); - -// Appends a primitive value to the list. -IREE_API_EXPORT iree_status_t IREE_API_CALL iree_vm_variant_list_append_value( - iree_vm_variant_list_t* list, iree_vm_value_t value); - -// Appends a ref object to the list by retaining it. -IREE_API_EXPORT iree_status_t IREE_API_CALL -iree_vm_variant_list_append_ref_retain(iree_vm_variant_list_t* list, - iree_vm_ref_t* ref); - -// Appends a ref object to the list by moving it. -IREE_API_EXPORT iree_status_t IREE_API_CALL -iree_vm_variant_list_append_ref_move(iree_vm_variant_list_t* list, - iree_vm_ref_t* ref); - -// Appends a null ref to the list. -IREE_API_EXPORT iree_status_t IREE_API_CALL -iree_vm_variant_list_append_null_ref(iree_vm_variant_list_t* list); - -// Returns a pointer to the variant element at the given index. -IREE_API_EXPORT iree_vm_variant_t* IREE_API_CALL -iree_vm_variant_list_get(iree_vm_variant_list_t* list, iree_host_size_t i); - -#endif // IREE_API_NO_PROTOTYPES - -#ifdef __cplusplus -} // extern "C" -#endif // __cplusplus - -#endif // IREE_VM_VARIANT_LIST_H_
diff --git a/kokoro/gcp_ubuntu/bazel/bindings/common.cfg b/kokoro/gcp_ubuntu/bazel/bindings/common.cfg index d4a4e26..8a49430 100644 --- a/kokoro/gcp_ubuntu/bazel/bindings/common.cfg +++ b/kokoro/gcp_ubuntu/bazel/bindings/common.cfg
@@ -17,4 +17,4 @@ # Common configuration for Kokoro builds that run the bindings build with bazel # on linux. -build_file: "iree/kokoro/gcp_ubuntu/bazel/bindings/build_kokoro.sh" +build_file: "iree/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/bindings/build_kokoro.sh"
diff --git a/kokoro/gcp_ubuntu/bazel/core/common.cfg b/kokoro/gcp_ubuntu/bazel/core/common.cfg index b8ccaa7..3a22d10 100755 --- a/kokoro/gcp_ubuntu/bazel/core/common.cfg +++ b/kokoro/gcp_ubuntu/bazel/core/common.cfg
@@ -17,4 +17,4 @@ # Common configuration for Kokoro builds that run the core build with bazel on # linux. -build_file: "iree/kokoro/gcp_ubuntu/bazel/core/build_kokoro.sh" +build_file: "iree/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/core/build_kokoro.sh"
diff --git a/kokoro/gcp_ubuntu/bazel/integrations/common.cfg b/kokoro/gcp_ubuntu/bazel/integrations/common.cfg index 9331508..eb31e55 100644 --- a/kokoro/gcp_ubuntu/bazel/integrations/common.cfg +++ b/kokoro/gcp_ubuntu/bazel/integrations/common.cfg
@@ -17,4 +17,4 @@ # Common configuration for Kokoro builds that run the integrations build with # bazel on linux. -build_file: "iree/kokoro/gcp_ubuntu/bazel/integrations/build_kokoro.sh" +build_file: "iree/build_tools/kokoro/gcp_ubuntu/bazel/linux/x86/integrations/build_kokoro.sh"
diff --git a/kokoro/gcp_ubuntu/cmake/android/arm64-v8a/common.cfg b/kokoro/gcp_ubuntu/cmake/android/arm64-v8a/common.cfg index 8159f36..1376e08 100644 --- a/kokoro/gcp_ubuntu/cmake/android/arm64-v8a/common.cfg +++ b/kokoro/gcp_ubuntu/cmake/android/arm64-v8a/common.cfg
@@ -17,4 +17,4 @@ # Common configuration for Kokoro builds that cross-compile IREE towards # Android arm64-v8a using CMake. -build_file: "iree/kokoro/gcp_ubuntu/cmake/android/arm64-v8a/build_kokoro.sh" +build_file: "iree/build_tools/kokoro/gcp_ubuntu/cmake/android/arm64-v8a/build_kokoro.sh"
diff --git a/kokoro/gcp_ubuntu/cmake/common.cfg b/kokoro/gcp_ubuntu/cmake/common.cfg index 838f4a5..49e6865 100644 --- a/kokoro/gcp_ubuntu/cmake/common.cfg +++ b/kokoro/gcp_ubuntu/cmake/common.cfg
@@ -16,4 +16,4 @@ # Common configuration for Kokoro builds that run cmake on linux. -build_file: "iree/kokoro/gcp_ubuntu/cmake/build_kokoro.sh" +build_file: "iree/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86/build_kokoro.sh"
diff --git a/kokoro/gcp_ubuntu/cmake/linux/x86-turing/common.cfg b/kokoro/gcp_ubuntu/cmake/linux/x86-turing/common.cfg index 9e6847d..bdb9163 100644 --- a/kokoro/gcp_ubuntu/cmake/linux/x86-turing/common.cfg +++ b/kokoro/gcp_ubuntu/cmake/linux/x86-turing/common.cfg
@@ -16,4 +16,4 @@ # Common configuration for Kokoro builds that run cmake on linux. -build_file: "iree/kokoro/gcp_ubuntu/cmake/linux/x86-turing/build_kokoro.sh" +build_file: "iree/build_tools/kokoro/gcp_ubuntu/cmake/linux/x86-turing/build_kokoro.sh"
diff --git a/packaging/python/common_setup.py b/packaging/python/common_setup.py index cbaf1d1..149dbdf 100644 --- a/packaging/python/common_setup.py +++ b/packaging/python/common_setup.py
@@ -16,6 +16,7 @@ import platform import setuptools import sys +import sysconfig from datetime import date @@ -99,15 +100,6 @@ } -def get_native_file_extension(): - if platform.system() == "Windows": - return "pyd" - elif platform.system() == "Darwin": - return "dylib" - else: - return "so" - - def setup(**kwargs): # See: https://stackoverflow.com/q/45150304 try: @@ -128,7 +120,7 @@ # Unfortunately, bazel is imprecise and scatters .so files around, so # need to be specific. package_data = { - "": ["*.%s" % (get_native_file_extension(),)], + "": ["*%s" % (sysconfig.get_config_var("EXT_SUFFIX"),)], } setuptools.setup( package_data=package_data,
diff --git a/packaging/python/setup_compiler.py b/packaging/python/setup_compiler.py index c33d7fd..ee9e87d 100644 --- a/packaging/python/setup_compiler.py +++ b/packaging/python/setup_compiler.py
@@ -39,7 +39,12 @@ print("Found packages:", packages) setup_kwargs = common_setup.get_setup_defaults( sub_project="compiler", description="IREE Generic Compiler") - common_setup.setup(packages=packages, **setup_kwargs) + common_setup.setup( + packages=packages, + ext_modules=[ + setuptools.Extension(name="pyiree.compiler.binding", sources=[]), + ], + **setup_kwargs) if __name__ == "__main__":
diff --git a/packaging/python/setup_rt.py b/packaging/python/setup_rt.py index ef39248..af30fd6 100644 --- a/packaging/python/setup_rt.py +++ b/packaging/python/setup_rt.py
@@ -35,7 +35,12 @@ setup_kwargs = common_setup.get_setup_defaults( sub_project="rt", description="IREE Runtime Components (for executing compiled programs)") - common_setup.setup(packages=packages, **setup_kwargs) + common_setup.setup( + packages=packages, + ext_modules=[ + setuptools.Extension(name="pyiree.rt.binding", sources=[]), + ], + **setup_kwargs) if __name__ == "__main__":
diff --git a/packaging/python/setup_tf.py b/packaging/python/setup_tf.py index 6f97d70..9eccf4f 100644 --- a/packaging/python/setup_tf.py +++ b/packaging/python/setup_tf.py
@@ -45,7 +45,12 @@ sub_project="tf", description="IREE TensorFlow Compiler", package_dir=package_dir) - common_setup.setup(packages=packages, **setup_kwargs) + common_setup.setup( + packages=packages, + ext_modules=[ + setuptools.Extension(name="pyiree.tf.compiler.binding", sources=[]), + ], + **setup_kwargs) if __name__ == "__main__":
diff --git a/scripts/prepare_doc_publication.py b/scripts/prepare_doc_publication.py index a4bbad9..cc396dc 100755 --- a/scripts/prepare_doc_publication.py +++ b/scripts/prepare_doc_publication.py
@@ -59,12 +59,8 @@ 'getting_started_android_cmake.md': 'Android with CMake', 'generic_vulkan_env_setup.md': 'Generic Vulkan Setup', 'getting_started_python.md': 'Python', - 'cmake_options_and_variables.md': 'CMake Options and Variables', - 'op_coverage.md': 'XLA HLO Operation Coverage', - 'e2e_coverage.md': 'TensorFlow E2E Coverage', - 'roadmap.md': 'Short-term Focus Areas', - 'roadmap_design.md': 'Long-term Design Roadmap', - 'iree_community.md': 'Community', + 'milestones.md': 'Short-term Focus Areas', + 'design_roadmap.md': 'Long-term Design Roadmap', } # A dictionary containing source file to permanent link mappings. @@ -75,25 +71,6 @@ # allows one to override the permanent link if necessary. PERMALINK_DICT = { 'index.md': '/', - 'getting_started_linux_bazel.md': 'GetStarted/LinuxBazel', - 'getting_started_linux_cmake.md': 'GetStarted/LinuxCMake', - 'getting_started_linux_vulkan.md': 'GetStarted/LinuxVulkan', - 'getting_started_windows_bazel.md': 'GetStarted/WindowsBazel', - 'getting_started_windows_cmake.md': 'GetStarted/WindowsCMake', - 'getting_started_windows_vulkan.md': 'GetStarted/WindowsVulkan', - 'getting_started_macos_cmake.md': 'GetStarted/macOSCMake', - 'getting_started_macos_vulkan.md': 'GetStarted/macOSVulkan', - 'getting_started_android_cmake.md': 'GetStarted/AndroidCMake', - 'generic_vulkan_env_setup.md': 'GetStarted/GenericVulkanSetup', - 'getting_started_python.md': 'GetStarted/Python', - 'cmake_options_and_variables.md': 'GetStarted/CMakeOptionsVariables', - 'developer_overview.md': 'DeveloperOverview', - 'testing_guide.md': 'TestingGuide', - 'op_coverage.md': 'HLOOpCoverage', - 'e2e_coverage.md': 'TensorFlowE2ECoverage', - 'roadmap.md': 'FocusAreas', - 'roadmap_design.md': 'DesignRoadmap', - 'iree_community.md': 'Community', } # A dictionary containing source file to navigation order mappings. @@ -102,15 +79,18 @@ # the left panel of https://google.github.io/iree website. This allows one # to specify an order for a specific doc. NAVI_ORDER_DICT = { + # Top level entries 'index.md': 1, - # 'Getting Started' is 2. - 'developer_overview.md': 3, - 'roadmap_design.md': 4, - 'roadmap.md': 5, - 'op_coverage.md': 6, - 'e2e_coverage.md': 7, - 'testing_guide.md': 8, + # 'Using IREE' is 2. + # 'Getting Started' is 3. + # 'Developing IREE' is 4. + 'design_roadmap.md': 5, + 'milestones.md': 6, + 'xla_op_coverage.md': 7, + 'tf_e2e_coverage.md': 8, 'iree_community.md': 9, + # 'Design Docs' is 10. + # 'Dialect Definitions' is 11. # Within 'Getting Started' use explicit ordering. # Alphabetical would put 'bazel' before 'cmake' and 'python' between 'linux' @@ -127,6 +107,16 @@ 'getting_started_python.md': 10, 'generic_vulkan_env_setup.md': 11, 'cmake_options_and_variables.md': 12, + + # Within 'Developing IREE' use explicit ordering. + 'developer_overview.md': 1, + 'contributor_tips.md': 2, + 'testing_guide.md': 3, + 'benchmarking.md': 4, + 'repository_management.md': 5, + + # Within 'Using IREE' use explicit ordering. + 'using_colab.md': 1, } # A dictionary containing source directory to section tile mappings. @@ -137,14 +127,17 @@ # Note that the title here must match with index.md file's title under the # subdirectory. DIRECTORY_TITLE_DICT = { + 'design_docs': 'Design Docs', + 'developing_iree': 'Developing IREE', 'Dialects': 'Dialect Definitions', - 'GetStarted': 'Getting Started', + 'get_started': 'Getting Started', + 'using_iree': 'Using IREE', } # A dictionary containing the supporting JavaScript files for each doc. JS_FILES_DICT = { - 'op_coverage.md': ['js/add_classes.js'], - 'e2e_coverage.md': ['js/add_classes.js'], + 'xla_op_coverage.md': ['js/add_classes.js'], + 'tf_e2e_coverage.md': ['js/add_classes.js'], } @@ -164,19 +157,20 @@ # Use the default layout for everything. front_matter['layout'] = 'default' # Use the base filename as permanent link. - front_matter['permalink'] = base_name + # Replace '_' with '-'. Underscores are not typical in URLs... + front_matter['permalink'] = base_name.replace('_', '-') # Organize each doc to a section matching its directory structure. if relpath and relpath != '.': - front_matter['parent'] = relpath - front_matter['permalink'] = f'{relpath}/{front_matter["permalink"]}' + hyphen_relpath = relpath.replace('_', '-') + front_matter['permalink'] = f'{hyphen_relpath}/{front_matter["permalink"]}' # Find the title and TOC. lines = content.splitlines() title_line_index = None toc_index = None for (index, line) in enumerate(lines): - if line.startswith('# '): + if line.startswith('# ') and title_line_index is None: title_line_index = index if line == '[TOC]': toc_index = index
diff --git a/scripts/update_e2e_coverage.py b/scripts/update_e2e_coverage.py index ea691e4..e9cf397 100755 --- a/scripts/update_e2e_coverage.py +++ b/scripts/update_e2e_coverage.py
@@ -64,6 +64,7 @@ - vulkan-spirv The table shows the supported TensorFlow functions and models on each backend. +It is auto-generated from IREE's test status. """ @@ -172,7 +173,7 @@ content.append(generate_table(test_suite)) content = '\n\n'.join(content) + '\n' # Trailing newline. - table_path = os.path.join(args.build_dir, 'doc', 'e2e_coverage.md') + table_path = os.path.join(args.build_dir, 'doc', 'tf_e2e_coverage.md') with open(table_path, 'w', encoding='utf-8') as f: f.write(E2E_COVERAGE_DESCRIPTION) f.write(content)
diff --git a/scripts/update_op_coverage.py b/scripts/update_op_coverage.py index b01b687..c2c6dd9 100755 --- a/scripts/update_op_coverage.py +++ b/scripts/update_op_coverage.py
@@ -30,14 +30,15 @@ E2E_XLA_OPS_PATH = 'iree/test/e2e/xla_ops' # TODO(scotttodd): LLVM AOT (dylib-llvm-aot) HAL target(s) -OP_COVERAGE_DESCRIPTION = """# HLO Op Coverage +OP_COVERAGE_DESCRIPTION = """# XLA HLO Op Coverage There are three backend [targets](https://github.com/google/iree/tree/main/iree/compiler/Dialect/HAL/Target) in IREE: - vmla - llvm-ir - vulkan-spirv -The table shows the supported XLA HLO ops on each backend. +The table shows the supported XLA HLO ops on each backend. It is auto-generated +from IREE's test status. """ @@ -117,7 +118,7 @@ if __name__ == '__main__': args = parse_arguments() content = generate_table(args.build_dir) - table_path = os.path.join(args.build_dir, 'doc', 'op_coverage.md') + table_path = os.path.join(args.build_dir, 'doc', 'xla_op_coverage.md') with open(table_path, 'w', encoding='utf-8') as f: f.write(OP_COVERAGE_DESCRIPTION) f.write(content)
diff --git a/third_party/flatcc b/third_party/flatcc new file mode 160000 index 0000000..4fb0ff7 --- /dev/null +++ b/third_party/flatcc
@@ -0,0 +1 @@ +Subproject commit 4fb0ff7069bd88ee85902f4d0bb62794e5f6d021
diff --git a/third_party/llvm-project b/third_party/llvm-project index de0c6bd..99ad956 160000 --- a/third_party/llvm-project +++ b/third_party/llvm-project
@@ -1 +1 @@ -Subproject commit de0c6bd56b41081f1b89a1c7a0bf2597fd6d0104 +Subproject commit 99ad956fdaee5398fdcf46fa49cb433cf52dc461
diff --git a/third_party/tensorflow b/third_party/tensorflow index e36aca0..e4a48da 160000 --- a/third_party/tensorflow +++ b/third_party/tensorflow
@@ -1 +1 @@ -Subproject commit e36aca0132fbcde0bc820d56185e3078f97a879d +Subproject commit e4a48da690fac3443825b535ec02cb31c5625337
