From a6043e2cb7907937d662c5e27d6642977b1e7acd Mon Sep 17 00:00:00 2001 From: Ben Vanik Date: Thu, 3 Oct 2024 08:49:55 -0700 Subject: [PATCH] Adding a null HAL driver. (#18675) This is a dummy that can be used as a template and documentation for new HAL drivers and a change detector for HAL APIs. It's also useful as a way to get the minimum size of a HAL driver with default emulation. --- CMakeLists.txt | 67 ++- runtime/src/iree/hal/drivers/BUILD.bazel | 6 + runtime/src/iree/hal/drivers/CMakeLists.txt | 12 +- runtime/src/iree/hal/drivers/init.c | 27 +- runtime/src/iree/hal/drivers/null/BUILD.bazel | 50 ++ .../src/iree/hal/drivers/null/CMakeLists.txt | 49 ++ runtime/src/iree/hal/drivers/null/README.md | 66 +++ runtime/src/iree/hal/drivers/null/allocator.c | 334 +++++++++++ runtime/src/iree/hal/drivers/null/allocator.h | 17 + runtime/src/iree/hal/drivers/null/api.h | 82 +++ runtime/src/iree/hal/drivers/null/buffer.c | 165 ++++++ runtime/src/iree/hal/drivers/null/buffer.h | 22 + runtime/src/iree/hal/drivers/null/channel.c | 127 ++++ runtime/src/iree/hal/drivers/null/channel.h | 18 + .../iree/hal/drivers/null/command_buffer.c | 367 ++++++++++++ .../iree/hal/drivers/null/command_buffer.h | 25 + runtime/src/iree/hal/drivers/null/device.c | 559 ++++++++++++++++++ runtime/src/iree/hal/drivers/null/device.h | 17 + runtime/src/iree/hal/drivers/null/driver.c | 201 +++++++ runtime/src/iree/hal/drivers/null/driver.h | 17 + runtime/src/iree/hal/drivers/null/event.c | 60 ++ runtime/src/iree/hal/drivers/null/event.h | 18 + .../src/iree/hal/drivers/null/executable.c | 86 +++ .../src/iree/hal/drivers/null/executable.h | 21 + .../iree/hal/drivers/null/executable_cache.c | 101 ++++ .../iree/hal/drivers/null/executable_cache.h | 23 + .../hal/drivers/null/registration/BUILD.bazel | 27 + .../drivers/null/registration/CMakeLists.txt | 29 + .../drivers/null/registration/driver_module.c | 60 ++ .../drivers/null/registration/driver_module.h | 24 + runtime/src/iree/hal/drivers/null/semaphore.c | 172 ++++++ runtime/src/iree/hal/drivers/null/semaphore.h | 23 + 32 files changed, 2832 insertions(+), 40 deletions(-) create mode 100644 runtime/src/iree/hal/drivers/null/BUILD.bazel create mode 100644 runtime/src/iree/hal/drivers/null/CMakeLists.txt create mode 100644 runtime/src/iree/hal/drivers/null/README.md create mode 100644 runtime/src/iree/hal/drivers/null/allocator.c create mode 100644 runtime/src/iree/hal/drivers/null/allocator.h create mode 100644 runtime/src/iree/hal/drivers/null/api.h create mode 100644 runtime/src/iree/hal/drivers/null/buffer.c create mode 100644 runtime/src/iree/hal/drivers/null/buffer.h create mode 100644 runtime/src/iree/hal/drivers/null/channel.c create mode 100644 runtime/src/iree/hal/drivers/null/channel.h create mode 100644 runtime/src/iree/hal/drivers/null/command_buffer.c create mode 100644 runtime/src/iree/hal/drivers/null/command_buffer.h create mode 100644 runtime/src/iree/hal/drivers/null/device.c create mode 100644 runtime/src/iree/hal/drivers/null/device.h create mode 100644 runtime/src/iree/hal/drivers/null/driver.c create mode 100644 runtime/src/iree/hal/drivers/null/driver.h create mode 100644 runtime/src/iree/hal/drivers/null/event.c create mode 100644 runtime/src/iree/hal/drivers/null/event.h create mode 100644 runtime/src/iree/hal/drivers/null/executable.c create mode 100644 runtime/src/iree/hal/drivers/null/executable.h create mode 100644 runtime/src/iree/hal/drivers/null/executable_cache.c create mode 100644 runtime/src/iree/hal/drivers/null/executable_cache.h create mode 100644 runtime/src/iree/hal/drivers/null/registration/BUILD.bazel create mode 100644 runtime/src/iree/hal/drivers/null/registration/CMakeLists.txt create mode 100644 runtime/src/iree/hal/drivers/null/registration/driver_module.c create mode 100644 runtime/src/iree/hal/drivers/null/registration/driver_module.h create mode 100644 runtime/src/iree/hal/drivers/null/semaphore.c create mode 100644 runtime/src/iree/hal/drivers/null/semaphore.h diff --git a/CMakeLists.txt b/CMakeLists.txt index 88a6d23baad7..571231ba34c7 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -257,13 +257,6 @@ set(IREE_HAL_DRIVER_CUDA_DEFAULT OFF) # required to actually run HIP workloads. set(IREE_HAL_DRIVER_HIP_DEFAULT OFF) -# Vulkan support is enabled by default if the platform might support Vulkan. -# Apple platforms support Metal instead of Vulkan, though MoltenVK may work. -set(IREE_HAL_DRIVER_VULKAN_DEFAULT ${IREE_HAL_DRIVER_DEFAULTS}) -if(APPLE) - set(IREE_HAL_DRIVER_VULKAN_DEFAULT OFF) -endif() - # Metal support is enabled if it's one of the Apple platforms. set(IREE_HAL_DRIVER_METAL_DEFAULT ${IREE_HAL_DRIVER_DEFAULTS}) # Right now only support Apple silicon devices. @@ -271,12 +264,29 @@ if(NOT APPLE OR NOT ${CMAKE_SYSTEM_PROCESSOR} MATCHES "arm64") set(IREE_HAL_DRIVER_METAL_DEFAULT OFF) endif() +# Null skeleton driver is only enabled in debug builds or dev mode. +# We don't want to ship release builds with it or count it when calculating +# binary sizes of minified builds. +set(IREE_HAL_DRIVER_NULL_DEFAULT OFF) +string(TOUPPER "${CMAKE_BUILD_TYPE}" _UPPERCASE_CMAKE_BUILD_TYPE) +if (IREE_DEV_MODE OR (_UPPERCASE_CMAKE_BUILD_TYPE STREQUAL "DEBUG")) + set(IREE_HAL_DRIVER_NULL_DEFAULT ON) +endif() + +# Vulkan support is enabled by default if the platform might support Vulkan. +# Apple platforms support Metal instead of Vulkan, though MoltenVK may work. +set(IREE_HAL_DRIVER_VULKAN_DEFAULT ${IREE_HAL_DRIVER_DEFAULTS}) +if(APPLE) + set(IREE_HAL_DRIVER_VULKAN_DEFAULT OFF) +endif() + option(IREE_HAL_DRIVER_CUDA "Enables the 'cuda' runtime HAL driver" ${IREE_HAL_DRIVER_CUDA_DEFAULT}) option(IREE_HAL_DRIVER_HIP "Enables the 'hip' runtime HAL driver" ${IREE_HAL_DRIVER_HIP_DEFAULT}) option(IREE_HAL_DRIVER_LOCAL_SYNC "Enables the 'local-sync' runtime HAL driver" ${IREE_HAL_DRIVER_DEFAULTS}) option(IREE_HAL_DRIVER_LOCAL_TASK "Enables the 'local-task' runtime HAL driver" ${IREE_HAL_DRIVER_DEFAULTS}) -option(IREE_HAL_DRIVER_VULKAN "Enables the 'vulkan' runtime HAL driver" ${IREE_HAL_DRIVER_VULKAN_DEFAULT}) option(IREE_HAL_DRIVER_METAL "Enables the 'metal' runtime HAL driver" ${IREE_HAL_DRIVER_METAL_DEFAULT}) +option(IREE_HAL_DRIVER_NULL "Enables the 'null' runtime HAL driver" ${IREE_HAL_DRIVER_NULL_DEFAULT}) +option(IREE_HAL_DRIVER_VULKAN "Enables the 'vulkan' runtime HAL driver" ${IREE_HAL_DRIVER_VULKAN_DEFAULT}) option(IREE_HAL_EXECUTABLE_LOADER_DEFAULTS "Sets the default value for all runtime HAL executable loaders" ON) set(IREE_HAL_EXECUTABLE_LOADER_EMBEDDED_ELF_DEFAULT ${IREE_HAL_EXECUTABLE_LOADER_DEFAULTS}) @@ -332,12 +342,15 @@ endif() if(IREE_HAL_DRIVER_LOCAL_TASK) message(STATUS " - local-task") endif() -if(IREE_HAL_DRIVER_VULKAN) - message(STATUS " - vulkan") -endif() if(IREE_HAL_DRIVER_METAL) message(STATUS " - metal") endif() +if(IREE_HAL_DRIVER_NULL) + message(STATUS " - null") +endif() +if(IREE_HAL_DRIVER_VULKAN) + message(STATUS " - vulkan") +endif() if(IREE_EXTERNAL_HAL_DRIVERS) message(STATUS " + external: ${IREE_EXTERNAL_HAL_DRIVERS}") endif() @@ -501,22 +514,22 @@ option(IREE_ENABLE_ASSERTIONS "Force unset of NDEBUG compile option" OFF) function(iree_fix_ndebug) string(TOUPPER "${CMAKE_BUILD_TYPE}" _UPPERCASE_CMAKE_BUILD_TYPE) if(IREE_ENABLE_ASSERTIONS AND NOT "${_UPPERCASE_CMAKE_BUILD_TYPE}" STREQUAL "DEBUG") - # Also remove /D NDEBUG to avoid MSVC warnings about conflicting defines. - foreach(_FLAGS_VAR_TO_SCRUB - CMAKE_CXX_FLAGS_${_UPPERCASE_CMAKE_BUILD_TYPE} - CMAKE_C_FLAGS_${_UPPERCASE_CMAKE_BUILD_TYPE}) - set(_ORIGINAL_FLAGS "${${_FLAGS_VAR_TO_SCRUB}}") - string(REGEX REPLACE "(^| )[/-]D *NDEBUG($| )" " " _ALTERED_FLAGS "${_ORIGINAL_FLAGS}") - if(NOT "${_ORIGINAL_FLAGS}" STREQUAL "${_ALTERED_FLAGS}") - message(STATUS - "IREE_ENABLE_ASSERTIONS force disabled NDEBUG for ${_FLAGS_VAR_TO_SCRUB}: '${_ORIGINAL_FLAGS}' -> '${_ALTERED_FLAGS}'") - set(${_FLAGS_VAR_TO_SCRUB} "${_ALTERED_FLAGS}" PARENT_SCOPE) - endif() - endforeach() - - # Make sure that LLVM doesn't add its own logic for assertion disabling. - # We'd like to make sure that we are not dueling over globals. - set(LLVM_ENABLE_ASSERTIONS OFF PARENT_SCOPE) + # Also remove /D NDEBUG to avoid MSVC warnings about conflicting defines. + foreach(_FLAGS_VAR_TO_SCRUB + CMAKE_CXX_FLAGS_${_UPPERCASE_CMAKE_BUILD_TYPE} + CMAKE_C_FLAGS_${_UPPERCASE_CMAKE_BUILD_TYPE}) + set(_ORIGINAL_FLAGS "${${_FLAGS_VAR_TO_SCRUB}}") + string(REGEX REPLACE "(^| )[/-]D *NDEBUG($| )" " " _ALTERED_FLAGS "${_ORIGINAL_FLAGS}") + if(NOT "${_ORIGINAL_FLAGS}" STREQUAL "${_ALTERED_FLAGS}") + message(STATUS + "IREE_ENABLE_ASSERTIONS force disabled NDEBUG for ${_FLAGS_VAR_TO_SCRUB}: '${_ORIGINAL_FLAGS}' -> '${_ALTERED_FLAGS}'") + set(${_FLAGS_VAR_TO_SCRUB} "${_ALTERED_FLAGS}" PARENT_SCOPE) + endif() + endforeach() + + # Make sure that LLVM doesn't add its own logic for assertion disabling. + # We'd like to make sure that we are not dueling over globals. + set(LLVM_ENABLE_ASSERTIONS OFF PARENT_SCOPE) endif() endfunction() iree_fix_ndebug() diff --git a/runtime/src/iree/hal/drivers/BUILD.bazel b/runtime/src/iree/hal/drivers/BUILD.bazel index f7b05ddf7d83..06a2a93fbea4 100644 --- a/runtime/src/iree/hal/drivers/BUILD.bazel +++ b/runtime/src/iree/hal/drivers/BUILD.bazel @@ -19,6 +19,7 @@ string_list_flag( "cuda", "local-sync", "local-task", + "null", "vulkan", ], ) @@ -26,6 +27,7 @@ string_list_flag( UNCONDITIONAL_DRIVERS = [ "local-sync", "local-task", + "null", "vulkan", ] @@ -67,6 +69,10 @@ iree_runtime_cc_library( ":local-task_enabled": ["//runtime/src/iree/hal/drivers/local_task/registration"], "//conditions:default": [], }) + + select({ + ":null_enabled": ["//runtime/src/iree/hal/drivers/null/registration"], + "//conditions:default": [], + }) + select({ ":vulkan_enabled": ["//runtime/src/iree/hal/drivers/vulkan/registration"], "//conditions:default": [], diff --git a/runtime/src/iree/hal/drivers/CMakeLists.txt b/runtime/src/iree/hal/drivers/CMakeLists.txt index b193b9994506..65e6ae712c54 100644 --- a/runtime/src/iree/hal/drivers/CMakeLists.txt +++ b/runtime/src/iree/hal/drivers/CMakeLists.txt @@ -140,14 +140,18 @@ if(IREE_HAL_DRIVER_LOCAL_TASK) add_subdirectory(local_task) list(APPEND _INIT_INTERNAL_DEPS iree::hal::drivers::local_task::registration) endif() -if(IREE_HAL_DRIVER_VULKAN) - add_subdirectory(vulkan) - list(APPEND _INIT_INTERNAL_DEPS iree::hal::drivers::vulkan::registration) -endif() if(IREE_HAL_DRIVER_METAL) add_subdirectory(metal) list(APPEND _INIT_INTERNAL_DEPS iree::hal::drivers::metal::registration) endif() +if(IREE_HAL_DRIVER_NULL) + add_subdirectory(null) + list(APPEND _INIT_INTERNAL_DEPS iree::hal::drivers::null::registration) +endif() +if(IREE_HAL_DRIVER_VULKAN) + add_subdirectory(vulkan) + list(APPEND _INIT_INTERNAL_DEPS iree::hal::drivers::vulkan::registration) +endif() iree_cc_library( NAME diff --git a/runtime/src/iree/hal/drivers/init.c b/runtime/src/iree/hal/drivers/init.c index a332ae5edc2d..57df102edbd1 100644 --- a/runtime/src/iree/hal/drivers/init.c +++ b/runtime/src/iree/hal/drivers/init.c @@ -20,16 +20,20 @@ #if defined(IREE_HAVE_HAL_LOCAL_TASK_DRIVER_MODULE) #include "iree/hal/drivers/local_task/registration/driver_module.h" + +#if defined(IREE_HAVE_HAL_METAL_DRIVER_MODULE) +#include "iree/hal/drivers/metal/registration/driver_module.h" +#endif // IREE_HAVE_HAL_METAL_DRIVER_MODULE #endif // IREE_HAVE_HAL_LOCAL_TASK_DRIVER_MODULE +#if defined(IREE_HAVE_HAL_NULL_DRIVER_MODULE) +#include "iree/hal/drivers/null/registration/driver_module.h" +#endif // IREE_HAVE_HAL_NULL_DRIVER_MODULE + #if defined(IREE_HAVE_HAL_VULKAN_DRIVER_MODULE) #include "iree/hal/drivers/vulkan/registration/driver_module.h" #endif // IREE_HAVE_HAL_VULKAN_DRIVER_MODULE -#if defined(IREE_HAVE_HAL_METAL_DRIVER_MODULE) -#include "iree/hal/drivers/metal/registration/driver_module.h" -#endif // IREE_HAVE_HAL_METAL_DRIVER_MODULE - #if defined(IREE_HAVE_HAL_EXTERNAL_DRIVERS) // Defined in the generated init_external.c file: extern iree_status_t iree_hal_register_external_drivers( @@ -65,16 +69,21 @@ iree_hal_register_all_available_drivers(iree_hal_driver_registry_t* registry) { z0, iree_hal_local_task_driver_module_register(registry)); #endif // IREE_HAVE_HAL_LOCAL_TASK_DRIVER_MODULE -#if defined(IREE_HAVE_HAL_VULKAN_DRIVER_MODULE) - IREE_RETURN_AND_END_ZONE_IF_ERROR( - z0, iree_hal_vulkan_driver_module_register(registry)); -#endif // IREE_HAVE_HAL_VULKAN_DRIVER_MODULE - #if defined(IREE_HAVE_HAL_METAL_DRIVER_MODULE) IREE_RETURN_AND_END_ZONE_IF_ERROR( z0, iree_hal_metal_driver_module_register(registry)); #endif // IREE_HAVE_HAL_METAL_DRIVER_MODULE +#if defined(IREE_HAVE_HAL_NULL_DRIVER_MODULE) + IREE_RETURN_AND_END_ZONE_IF_ERROR( + z0, iree_hal_null_driver_module_register(registry)); +#endif // IREE_HAVE_HAL_NULL_DRIVER_MODULE + +#if defined(IREE_HAVE_HAL_VULKAN_DRIVER_MODULE) + IREE_RETURN_AND_END_ZONE_IF_ERROR( + z0, iree_hal_vulkan_driver_module_register(registry)); +#endif // IREE_HAVE_HAL_VULKAN_DRIVER_MODULE + IREE_RETURN_AND_END_ZONE_IF_ERROR( z0, iree_hal_register_external_drivers(registry)); diff --git a/runtime/src/iree/hal/drivers/null/BUILD.bazel b/runtime/src/iree/hal/drivers/null/BUILD.bazel new file mode 100644 index 000000000000..a33a99c31c8c --- /dev/null +++ b/runtime/src/iree/hal/drivers/null/BUILD.bazel @@ -0,0 +1,50 @@ +# Copyright 2024 The IREE Authors +# +# Licensed under the Apache License v2.0 with LLVM Exceptions. +# See https://llvm.org/LICENSE.txt for license information. +# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +load("//build_tools/bazel:build_defs.oss.bzl", "iree_runtime_cc_library") + +package( + default_visibility = ["//visibility:public"], + features = ["layering_check"], + licenses = ["notice"], # Apache 2.0 +) + +iree_runtime_cc_library( + name = "null", + srcs = [ + "allocator.c", + "allocator.h", + "buffer.c", + "buffer.h", + "channel.c", + "channel.h", + "command_buffer.c", + "command_buffer.h", + "device.c", + "device.h", + "driver.c", + "driver.h", + "event.c", + "event.h", + "executable.c", + "executable.h", + "executable_cache.c", + "executable_cache.h", + "semaphore.c", + "semaphore.h", + ], + hdrs = [ + "api.h", + ], + deps = [ + "//runtime/src/iree/base", + "//runtime/src/iree/base/internal", + "//runtime/src/iree/hal", + "//runtime/src/iree/hal/utils:file_transfer", + "//runtime/src/iree/hal/utils:memory_file", + "//runtime/src/iree/hal/utils:semaphore_base", + ], +) diff --git a/runtime/src/iree/hal/drivers/null/CMakeLists.txt b/runtime/src/iree/hal/drivers/null/CMakeLists.txt new file mode 100644 index 000000000000..c5ed9eb49d67 --- /dev/null +++ b/runtime/src/iree/hal/drivers/null/CMakeLists.txt @@ -0,0 +1,49 @@ +################################################################################ +# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from # +# runtime/src/iree/hal/drivers/null/BUILD.bazel # +# # +# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary # +# CMake-only content. # +# # +# To disable autogeneration for this file entirely, delete this header. # +################################################################################ + +iree_add_all_subdirs() + +iree_cc_library( + NAME + null + HDRS + "api.h" + SRCS + "allocator.c" + "allocator.h" + "buffer.c" + "buffer.h" + "channel.c" + "channel.h" + "command_buffer.c" + "command_buffer.h" + "device.c" + "device.h" + "driver.c" + "driver.h" + "event.c" + "event.h" + "executable.c" + "executable.h" + "executable_cache.c" + "executable_cache.h" + "semaphore.c" + "semaphore.h" + DEPS + iree::base + iree::base::internal + iree::hal + iree::hal::utils::file_transfer + iree::hal::utils::memory_file + iree::hal::utils::semaphore_base + PUBLIC +) + +### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ### diff --git a/runtime/src/iree/hal/drivers/null/README.md b/runtime/src/iree/hal/drivers/null/README.md new file mode 100644 index 000000000000..3c3e4200334b --- /dev/null +++ b/runtime/src/iree/hal/drivers/null/README.md @@ -0,0 +1,66 @@ +# Null HAL Driver (`null`) + +This is a skeleton HAL driver that implements most interfaces needed by a HAL +driver. It contains notes that can be useful to implementers, stubs for +easy copy/paste/renaming, and some default behavior in places that many +implementations can use (such as file queue operations). It doesn't run anything +and nearly all methods return `UNIMPLEMENTED` errors: this can be used to +incrementally build a driver by running until you hit an UNIMPLEMENTED, +implementing it, and running again. Note however that the HAL is fairly regular +and starting with the files and running down the methods is often much easier +to do than the trial-and-error approach: if you can implement command buffer +fill (memset) you can often implement copy (memcpy) as well at the same time. + +## Instructions for Cloning + +1. Duplicate the entire directory in your own repository or the IREE + `experimental/` folder if going in-tree. +1. Find/replace `{Null}` with the friendly name of your driver (e.g. `Vulkan`). +1. Find/replace `_null_` with the C name of your driver (e.g. `vulkan`). +1. Find/replace `// TODO(null):` with your github ID, your driver name, or a + GitHub issue number tracking driver creation (e.g. `// TODO(#1234):`). + +## Build Setup + +HAL drivers are setup by adding some specially named cmake variables and then +pointing the IREE build at them by name. Projects embedding IREE runtime builds +as a submodule can use the `iree_register_external_hal_driver` cmake function +to do this, set the variables on the command line during cmake configure, or +via top-level project `CMakeLists.txt` before adding the IREE subdirectory. + +See [runtime/src/iree/hal/drivers/CMakeLists.txt](runtime/src/iree/hal/drivers/CMakeLists.txt) for more information. + +Example using the helper function: +```cmake +iree_register_external_hal_driver( + NAME + webgpu + SOURCE_DIR + "${CMAKE_CURRENT_SOURCE_DIR}/experimental/webgpu" + BINARY_DIR + "${CMAKE_CURRENT_BINARY_DIR}/experimental/webgpu" + DRIVER_TARGET + iree::experimental::webgpu::registration + REGISTER_FN + iree_hal_webgpu_driver_module_register +) +set(IREE_EXTERNAL_HAL_DRIVERS my_driver) +``` + +Example using the command line: +```sh +cmake ... \ + -DIREE_EXTERNAL_MY_DRIVER_HAL_DRIVER_TARGET=my_driver_static_library \ + -DIREE_EXTERNAL_MY_DRIVER_HAL_DRIVER_REGISTER=my_driver_register \ + -DIREE_EXTERNAL_HAL_DRIVERS=my_driver +``` + +## In-tree Drivers (`iree/hal/drivers/...`) + +IREE is generally conservative about hosting in-tree HAL drivers unless authored +by the core team or an SLA is in-place and maintained. Any new HAL drivers +should expect to start in forks or external repositories and not be expected +to merge without deep involvement with the IREE maintainers. IREE is not a +monorepo and it's perfectly fine to be out-of-tree. If ergonomics issues are +encountered with being out of tree please file issues so that support can be +improved. diff --git a/runtime/src/iree/hal/drivers/null/allocator.c b/runtime/src/iree/hal/drivers/null/allocator.c new file mode 100644 index 000000000000..f84f00257ee5 --- /dev/null +++ b/runtime/src/iree/hal/drivers/null/allocator.c @@ -0,0 +1,334 @@ +// Copyright 2024 The IREE Authors +// +// Licensed under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#include "iree/hal/drivers/null/allocator.h" + +#include "iree/hal/drivers/null/buffer.h" + +// TODO(null): use one ID per address space or pool - each shows as a different +// track in tracing tools. +#if IREE_TRACING_FEATURES & IREE_TRACING_FEATURE_ALLOCATION_TRACKING +static const char* IREE_HAL_NULL_ALLOCATOR_ID = "{Null} unpooled"; +#endif // IREE_TRACING_FEATURE_ALLOCATION_TRACKING + +typedef struct iree_hal_null_allocator_t { + iree_hal_resource_t resource; + iree_allocator_t host_allocator; + + IREE_STATISTICS(iree_hal_allocator_statistics_t statistics;) +} iree_hal_null_allocator_t; + +static const iree_hal_allocator_vtable_t iree_hal_null_allocator_vtable; + +static iree_hal_null_allocator_t* iree_hal_null_allocator_cast( + iree_hal_allocator_t* base_value) { + IREE_HAL_ASSERT_TYPE(base_value, &iree_hal_null_allocator_vtable); + return (iree_hal_null_allocator_t*)base_value; +} + +iree_status_t iree_hal_null_allocator_create( + iree_allocator_t host_allocator, iree_hal_allocator_t** out_allocator) { + IREE_ASSERT_ARGUMENT(out_allocator); + IREE_TRACE_ZONE_BEGIN(z0); + + iree_hal_null_allocator_t* allocator = NULL; + IREE_RETURN_AND_END_ZONE_IF_ERROR( + z0, iree_allocator_malloc(host_allocator, sizeof(*allocator), + (void**)&allocator)); + iree_hal_resource_initialize(&iree_hal_null_allocator_vtable, + &allocator->resource); + allocator->host_allocator = host_allocator; + + // TODO(null): query device heaps, supported features (concurrent access/etc), + // and prepare any pools that will be used during allocation. It's expected + // that most failures that occur after creation are allocation + // request-specific so preparing here will help keep the errors more + // localized. + iree_status_t status = + iree_make_status(IREE_STATUS_UNIMPLEMENTED, "allocator not implemented"); + + if (iree_status_is_ok(status)) { + *out_allocator = (iree_hal_allocator_t*)allocator; + } else { + iree_hal_allocator_release((iree_hal_allocator_t*)allocator); + } + IREE_TRACE_ZONE_END(z0); + return status; +} + +static void iree_hal_null_allocator_destroy( + iree_hal_allocator_t* IREE_RESTRICT base_allocator) { + IREE_ASSERT_ARGUMENT(base_allocator); + iree_hal_null_allocator_t* allocator = + iree_hal_null_allocator_cast(base_allocator); + IREE_TRACE_ZONE_BEGIN(z0); + + iree_allocator_free(allocator->host_allocator, allocator); + + IREE_TRACE_ZONE_END(z0); +} + +static iree_allocator_t iree_hal_null_allocator_host_allocator( + const iree_hal_allocator_t* IREE_RESTRICT base_allocator) { + iree_hal_null_allocator_t* allocator = + (iree_hal_null_allocator_t*)base_allocator; + return allocator->host_allocator; +} + +static iree_status_t iree_hal_null_allocator_trim( + iree_hal_allocator_t* IREE_RESTRICT base_allocator) { + iree_hal_null_allocator_t* allocator = + (iree_hal_null_allocator_t*)base_allocator; + + // TODO(null): if the allocator is retaining any unused resources they should + // be dropped here. If the underlying implementation has pools or caches it + // should be notified that a trim is requested. This is called in low-memory + // situations or when IREE is not going to be used for awhile (low power modes + // or suspension). + (void)allocator; + + return iree_ok_status(); +} + +static void iree_hal_null_allocator_query_statistics( + iree_hal_allocator_t* IREE_RESTRICT base_allocator, + iree_hal_allocator_statistics_t* IREE_RESTRICT out_statistics) { + IREE_STATISTICS({ + iree_hal_null_allocator_t* allocator = + iree_hal_null_allocator_cast(base_allocator); + memcpy(out_statistics, &allocator->statistics, sizeof(*out_statistics)); + // TODO(null): update statistics (merge). + }); +} + +static iree_status_t iree_hal_null_allocator_query_memory_heaps( + iree_hal_allocator_t* IREE_RESTRICT base_allocator, + iree_host_size_t capacity, + iree_hal_allocator_memory_heap_t* IREE_RESTRICT heaps, + iree_host_size_t* IREE_RESTRICT out_count) { + iree_hal_null_allocator_t* allocator = + iree_hal_null_allocator_cast(base_allocator); + + // TODO(null): return heap information. This is called at least once with a + // capacity that may be 0 (indicating a query for the total count) and the + // heaps should only be populated if capacity is sufficient to store all of + // them. + (void)allocator; + iree_status_t status = + iree_make_status(IREE_STATUS_UNIMPLEMENTED, "heap query not implemented"); + + return status; +} + +static iree_hal_buffer_compatibility_t +iree_hal_null_allocator_query_buffer_compatibility( + iree_hal_allocator_t* IREE_RESTRICT base_allocator, + iree_hal_buffer_params_t* IREE_RESTRICT params, + iree_device_size_t* IREE_RESTRICT allocation_size) { + iree_hal_null_allocator_t* allocator = + iree_hal_null_allocator_cast(base_allocator); + + // TODO(null): set compatibility rules based on the implementation. + // Note that the user may have requested that the allocator place the + // allocation based on whatever is optimal for the indicated usage by + // including the IREE_HAL_MEMORY_TYPE_OPTIMAL flag. It's still required that + // the implementation meet all the requirements but it is free to place it in + // either host or device memory so long as the appropriate bits are updated to + // indicate where it landed. + (void)allocator; + iree_hal_buffer_compatibility_t compatibility = + IREE_HAL_BUFFER_COMPATIBILITY_NONE; + + // We are now optimal. + params->type &= ~IREE_HAL_MEMORY_TYPE_OPTIMAL; + + // Guard against the corner case where the requested buffer size is 0. The + // application is unlikely to do anything when requesting a 0-byte buffer; but + // it can happen in real world use cases. So we should at least not crash. + if (*allocation_size == 0) *allocation_size = 4; + + return compatibility; +} + +static iree_status_t iree_hal_null_allocator_allocate_buffer( + iree_hal_allocator_t* IREE_RESTRICT base_allocator, + const iree_hal_buffer_params_t* IREE_RESTRICT params, + iree_device_size_t allocation_size, + iree_hal_buffer_t** IREE_RESTRICT out_buffer) { + iree_hal_null_allocator_t* allocator = + iree_hal_null_allocator_cast(base_allocator); + + // Coerce options into those required by the current device. + iree_hal_buffer_params_t compat_params = *params; + iree_hal_buffer_compatibility_t compatibility = + iree_hal_null_allocator_query_buffer_compatibility( + base_allocator, &compat_params, &allocation_size); + if (!iree_all_bits_set(compatibility, + IREE_HAL_BUFFER_COMPATIBILITY_ALLOCATABLE)) { + // TODO(benvanik): make a helper for this. +#if IREE_STATUS_MODE + iree_bitfield_string_temp_t temp0, temp1, temp2; + iree_string_view_t memory_type_str = + iree_hal_memory_type_format(params->type, &temp0); + iree_string_view_t usage_str = + iree_hal_buffer_usage_format(params->usage, &temp1); + iree_string_view_t compatibility_str = + iree_hal_buffer_compatibility_format(compatibility, &temp2); + return iree_make_status( + IREE_STATUS_INVALID_ARGUMENT, + "allocator cannot allocate a buffer with the given parameters; " + "memory_type=%.*s, usage=%.*s, compatibility=%.*s", + (int)memory_type_str.size, memory_type_str.data, (int)usage_str.size, + usage_str.data, (int)compatibility_str.size, compatibility_str.data); +#else + return iree_make_status( + IREE_STATUS_INVALID_ARGUMENT, + "allocator cannot allocate a buffer with the given parameters"); +#endif // IREE_STATUS_MODE + } + + // TODO(null): allocate the underlying device memory. The impl_ptr is just + // used for accounting and can be an opaque value (handle/etc) so long as it + // is consistent between the alloc and free and unique to the buffer while it + // is live. An example iree_hal_null_buffer_wrap is provided that can be used + // for implementations that are managing memory using underlying allocators + // and just wrapping those device pointers in the HAL buffer type. Other + // implementations that require more tracking can provide their own buffer + // types that do such tracking for them. + (void)allocator; + void* impl_ptr = NULL; + (void)impl_ptr; + iree_hal_buffer_t* buffer = NULL; + iree_status_t status = iree_make_status(IREE_STATUS_UNIMPLEMENTED, + "buffer allocation not implemented"); + + if (iree_status_is_ok(status)) { + // TODO(null): ensure this accounting is balanced in deallocate_buffer. + IREE_TRACE_ALLOC_NAMED(IREE_HAL_NULL_ALLOCATOR_ID, impl_ptr, + allocation_size); + IREE_STATISTICS(iree_hal_allocator_statistics_record_alloc( + &allocator->statistics, compat_params.type, allocation_size)); + *out_buffer = buffer; + } else { + iree_hal_buffer_release(buffer); + } + return status; +} + +static void iree_hal_null_allocator_deallocate_buffer( + iree_hal_allocator_t* IREE_RESTRICT base_allocator, + iree_hal_buffer_t* IREE_RESTRICT base_buffer) { + iree_hal_null_allocator_t* allocator = + iree_hal_null_allocator_cast(base_allocator); + + // TODO(null): free the underlying device memory here. Buffers allocated from + // this allocator will call this method to handle cleanup. Note that because + // this method is responsible for doing the base iree_hal_buffer_destroy and + // the caller assumes the memory has been freed an implementation could pool + // the buffer handle and return it in the future. + (void)allocator; + void* impl_ptr = NULL; + (void)impl_ptr; + + // TODO(null): if the buffer was imported then this accounting may need to be + // conditional depending on the implementation. + bool was_imported = false; + if (!was_imported) { + IREE_TRACE_FREE_NAMED(IREE_HAL_NULL_ALLOCATOR_ID, impl_ptr); + IREE_STATISTICS(iree_hal_allocator_statistics_record_free( + &allocator->statistics, iree_hal_buffer_memory_type(base_buffer), + iree_hal_buffer_allocation_size(base_buffer))); + } + + iree_hal_buffer_destroy(base_buffer); +} + +static iree_status_t iree_hal_null_allocator_import_buffer( + iree_hal_allocator_t* IREE_RESTRICT base_allocator, + const iree_hal_buffer_params_t* IREE_RESTRICT params, + iree_hal_external_buffer_t* IREE_RESTRICT external_buffer, + iree_hal_buffer_release_callback_t release_callback, + iree_hal_buffer_t** IREE_RESTRICT out_buffer) { + iree_hal_null_allocator_t* allocator = + iree_hal_null_allocator_cast(base_allocator); + + // Coerce options into those required by the current device. + iree_hal_buffer_params_t compat_params = *params; + iree_device_size_t allocation_size = external_buffer->size; + iree_hal_buffer_compatibility_t compatibility = + iree_hal_null_allocator_query_buffer_compatibility( + base_allocator, &compat_params, &allocation_size); + if (!iree_all_bits_set(compatibility, + IREE_HAL_BUFFER_COMPATIBILITY_IMPORTABLE)) { + // TODO(benvanik): make a helper for this. +#if IREE_STATUS_MODE + iree_bitfield_string_temp_t temp0, temp1, temp2; + iree_string_view_t memory_type_str = + iree_hal_memory_type_format(params->type, &temp0); + iree_string_view_t usage_str = + iree_hal_buffer_usage_format(params->usage, &temp1); + iree_string_view_t compatibility_str = + iree_hal_buffer_compatibility_format(compatibility, &temp2); + return iree_make_status( + IREE_STATUS_INVALID_ARGUMENT, + "allocator cannot import a buffer with the given parameters; " + "memory_type=%.*s, usage=%.*s, compatibility=%.*s", + (int)memory_type_str.size, memory_type_str.data, (int)usage_str.size, + usage_str.data, (int)compatibility_str.size, compatibility_str.data); +#else + return iree_make_status( + IREE_STATUS_INVALID_ARGUMENT, + "allocator cannot import a buffer with the given parameters"); +#endif // IREE_STATUS_MODE + } + + // TODO(null): switch on external_buffer->type and import the buffer. See the + // headers for more information on semantics. Most implementations can service + // IREE_HAL_EXTERNAL_BUFFER_TYPE_DEVICE_ALLOCATION by just wrapping the + // underlying device pointer. Those that can service + // IREE_HAL_EXTERNAL_BUFFER_TYPE_HOST_ALLOCATION may be able to avoid a lot of + // additional copies when moving data around between host and device or across + // devices from different drivers. + (void)allocator; + iree_status_t status = iree_make_status(IREE_STATUS_UNIMPLEMENTED, + "external buffer type not supported"); + + return status; +} + +static iree_status_t iree_hal_null_allocator_export_buffer( + iree_hal_allocator_t* IREE_RESTRICT base_allocator, + iree_hal_buffer_t* IREE_RESTRICT buffer, + iree_hal_external_buffer_type_t requested_type, + iree_hal_external_buffer_flags_t requested_flags, + iree_hal_external_buffer_t* IREE_RESTRICT out_external_buffer) { + iree_hal_null_allocator_t* allocator = + iree_hal_null_allocator_cast(base_allocator); + + // TODO(null): switch on requested_type and export as appropriate. Most + // implementations can service IREE_HAL_EXTERNAL_BUFFER_TYPE_DEVICE_ALLOCATION + // by just exposing the underlying device pointer. Those that can service + // IREE_HAL_EXTERNAL_BUFFER_TYPE_HOST_ALLOCATION may be able to avoid a lot of + // additional copies when moving data around between host and device or across + // devices from different drivers. + (void)allocator; + return iree_make_status(IREE_STATUS_UNAVAILABLE, + "external buffer type not supported"); +} + +static const iree_hal_allocator_vtable_t iree_hal_null_allocator_vtable = { + .destroy = iree_hal_null_allocator_destroy, + .host_allocator = iree_hal_null_allocator_host_allocator, + .trim = iree_hal_null_allocator_trim, + .query_statistics = iree_hal_null_allocator_query_statistics, + .query_memory_heaps = iree_hal_null_allocator_query_memory_heaps, + .query_buffer_compatibility = + iree_hal_null_allocator_query_buffer_compatibility, + .allocate_buffer = iree_hal_null_allocator_allocate_buffer, + .deallocate_buffer = iree_hal_null_allocator_deallocate_buffer, + .import_buffer = iree_hal_null_allocator_import_buffer, + .export_buffer = iree_hal_null_allocator_export_buffer, +}; diff --git a/runtime/src/iree/hal/drivers/null/allocator.h b/runtime/src/iree/hal/drivers/null/allocator.h new file mode 100644 index 000000000000..c0286bac6041 --- /dev/null +++ b/runtime/src/iree/hal/drivers/null/allocator.h @@ -0,0 +1,17 @@ +// Copyright 2024 The IREE Authors +// +// Licensed under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#ifndef IREE_HAL_DRIVERS_NULL_ALLOCATOR_H_ +#define IREE_HAL_DRIVERS_NULL_ALLOCATOR_H_ + +#include "iree/base/api.h" +#include "iree/hal/api.h" + +// Creates a {Null} buffer allocator used for persistent allocations. +iree_status_t iree_hal_null_allocator_create( + iree_allocator_t host_allocator, iree_hal_allocator_t** out_allocator); + +#endif // IREE_HAL_DRIVERS_NULL_ALLOCATOR_H_ diff --git a/runtime/src/iree/hal/drivers/null/api.h b/runtime/src/iree/hal/drivers/null/api.h new file mode 100644 index 000000000000..ad471735684e --- /dev/null +++ b/runtime/src/iree/hal/drivers/null/api.h @@ -0,0 +1,82 @@ +// Copyright 2024 The IREE Authors +// +// Licensed under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#ifndef IREE_HAL_DRIVERS_NULL_API_H_ +#define IREE_HAL_DRIVERS_NULL_API_H_ + +#include "iree/base/api.h" +#include "iree/hal/api.h" + +#ifdef __cplusplus +extern "C" { +#endif // __cplusplus + +//===----------------------------------------------------------------------===// +// iree_hal_null_device_t +//===----------------------------------------------------------------------===// + +// Parameters configuring an iree_hal_null_device_t. +// Must be initialized with iree_hal_null_device_options_initialize prior to +// use. +typedef struct iree_hal_null_device_options_t { + // TODO(null): options for initializing a device such as hardware identifiers, + // implementation mode switches, and debugging control. + int reserved; +} iree_hal_null_device_options_t; + +// Initializes |out_params| to default values. +IREE_API_EXPORT void iree_hal_null_device_options_initialize( + iree_hal_null_device_options_t* out_params); + +// Creates a {Null} HAL device with the given |params|. +// +// The provided |identifier| will be used by programs to distinguish the device +// type from other HAL implementations. If compiling programs with the IREE +// compiler this must match the value used by `IREE::HAL::TargetDevice`. +// +// |out_device| must be released by the caller (see iree_hal_device_release). +IREE_API_EXPORT iree_status_t iree_hal_null_device_create( + iree_string_view_t identifier, + const iree_hal_null_device_options_t* options, + iree_allocator_t host_allocator, iree_hal_device_t** out_device); + +//===----------------------------------------------------------------------===// +// iree_hal_null_driver_t +//===----------------------------------------------------------------------===// + +// Parameters for configuring an iree_hal_null_driver_t. +// Must be initialized with iree_hal_null_driver_options_initialize prior to +// use. +typedef struct iree_hal_null_driver_options_t { + // TODO(null): options for initializing the driver such as library search + // paths, version min/max, etc. + + // Default device options when none are provided during device creation. + iree_hal_null_device_options_t default_device_options; +} iree_hal_null_driver_options_t; + +// Initializes the given |out_options| with default driver creation options. +IREE_API_EXPORT void iree_hal_null_driver_options_initialize( + iree_hal_null_driver_options_t* out_options); + +// Creates a {Null} HAL driver with the given |options|, from which {Null} +// devices can be enumerated and created with specific parameters. +// +// The provided |identifier| will be used by programs to distinguish the device +// type from other HAL implementations. If compiling programs with the IREE +// compiler this must match the value used by IREE::HAL::TargetDevice. +// +// |out_driver| must be released by the caller (see iree_hal_driver_release). +IREE_API_EXPORT iree_status_t iree_hal_null_driver_create( + iree_string_view_t identifier, + const iree_hal_null_driver_options_t* options, + iree_allocator_t host_allocator, iree_hal_driver_t** out_driver); + +#ifdef __cplusplus +} // extern "C" +#endif // __cplusplus + +#endif // IREE_HAL_DRIVERS_NULL_API_H_ diff --git a/runtime/src/iree/hal/drivers/null/buffer.c b/runtime/src/iree/hal/drivers/null/buffer.c new file mode 100644 index 000000000000..f6eeecb11f20 --- /dev/null +++ b/runtime/src/iree/hal/drivers/null/buffer.c @@ -0,0 +1,165 @@ +// Copyright 2024 The IREE Authors +// +// Licensed under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#include "iree/hal/drivers/null/buffer.h" + +typedef struct iree_hal_null_buffer_t { + iree_hal_buffer_t base; + iree_hal_buffer_release_callback_t release_callback; +} iree_hal_null_buffer_t; + +static const iree_hal_buffer_vtable_t iree_hal_null_buffer_vtable; + +static iree_hal_null_buffer_t* iree_hal_null_buffer_cast( + iree_hal_buffer_t* base_value) { + IREE_HAL_ASSERT_TYPE(base_value, &iree_hal_null_buffer_vtable); + return (iree_hal_null_buffer_t*)base_value; +} + +static const iree_hal_null_buffer_t* iree_hal_null_buffer_const_cast( + const iree_hal_buffer_t* base_value) { + IREE_HAL_ASSERT_TYPE(base_value, &iree_hal_null_buffer_vtable); + return (const iree_hal_null_buffer_t*)base_value; +} + +iree_status_t iree_hal_null_buffer_wrap( + iree_hal_allocator_t* allocator, iree_hal_memory_type_t memory_type, + iree_hal_memory_access_t allowed_access, + iree_hal_buffer_usage_t allowed_usage, iree_device_size_t allocation_size, + iree_device_size_t byte_offset, iree_device_size_t byte_length, + iree_hal_buffer_release_callback_t release_callback, + iree_allocator_t host_allocator, iree_hal_buffer_t** out_buffer) { + IREE_ASSERT_ARGUMENT(out_buffer); + *out_buffer = NULL; + IREE_TRACE_ZONE_BEGIN(z0); + + iree_hal_null_buffer_t* buffer = NULL; + IREE_RETURN_AND_END_ZONE_IF_ERROR( + z0, + iree_allocator_malloc(host_allocator, sizeof(*buffer), (void**)&buffer)); + iree_hal_buffer_initialize(host_allocator, allocator, &buffer->base, + allocation_size, byte_offset, byte_length, + memory_type, allowed_access, allowed_usage, + &iree_hal_null_buffer_vtable, &buffer->base); + buffer->release_callback = release_callback; + + // TODO(null): retain or take ownership of provided handles/pointers/etc. + // Implementations may want to pass in an internal buffer type discriminator + // if there are multiple or use different top-level iree_hal_buffer_t + // implementations. + iree_status_t status = iree_make_status(IREE_STATUS_UNIMPLEMENTED, + "buffer wrapping not implemented"); + + if (iree_status_is_ok(status)) { + *out_buffer = &buffer->base; + } else { + iree_hal_buffer_release(&buffer->base); + } + IREE_TRACE_ZONE_END(z0); + return status; +} + +static void iree_hal_null_buffer_destroy(iree_hal_buffer_t* base_buffer) { + iree_hal_null_buffer_t* buffer = iree_hal_null_buffer_cast(base_buffer); + iree_allocator_t host_allocator = base_buffer->host_allocator; + IREE_TRACE_ZONE_BEGIN(z0); + + // Optionally call a release callback when the buffer is destroyed. Not all + // implementations may require this but it's cheap and provides additional + // flexibility. + if (buffer->release_callback.fn) { + buffer->release_callback.fn(buffer->release_callback.user_data, + base_buffer); + } + + iree_allocator_free(host_allocator, buffer); + + IREE_TRACE_ZONE_END(z0); +} + +static iree_status_t iree_hal_null_buffer_map_range( + iree_hal_buffer_t* base_buffer, iree_hal_mapping_mode_t mapping_mode, + iree_hal_memory_access_t memory_access, + iree_device_size_t local_byte_offset, iree_device_size_t local_byte_length, + iree_hal_buffer_mapping_t* mapping) { + iree_hal_null_buffer_t* buffer = iree_hal_null_buffer_cast(base_buffer); + + IREE_RETURN_IF_ERROR(iree_hal_buffer_validate_memory_type( + iree_hal_buffer_memory_type(base_buffer), + IREE_HAL_MEMORY_TYPE_HOST_VISIBLE)); + IREE_RETURN_IF_ERROR(iree_hal_buffer_validate_usage( + iree_hal_buffer_allowed_usage(base_buffer), + mapping_mode == IREE_HAL_MAPPING_MODE_PERSISTENT + ? IREE_HAL_BUFFER_USAGE_MAPPING_PERSISTENT + : IREE_HAL_BUFFER_USAGE_MAPPING_SCOPED)); + + // TODO(null): perform mapping as described. Note that local-to-buffer range + // adjustment may be required. The resulting mapping is populated with + // standard information such as contents indicating the host addressable + // memory range of the mapped buffer and implementation-specific information + // if additional resources are required. iree_hal_buffer_emulated_map_range + // can be used by implementations that have no way of providing host pointers + // at a large cost (alloc + device->host transfer on map and host->device + // transfer + dealloc on umap). Try not to use that. + (void)buffer; + iree_status_t status = iree_make_status(IREE_STATUS_UNIMPLEMENTED, + "buffer mapping not implemented"); + + return status; +} + +static iree_status_t iree_hal_null_buffer_unmap_range( + iree_hal_buffer_t* base_buffer, iree_device_size_t local_byte_offset, + iree_device_size_t local_byte_length, iree_hal_buffer_mapping_t* mapping) { + iree_hal_null_buffer_t* buffer = iree_hal_null_buffer_cast(base_buffer); + + // TODO(null): reverse of map_range. Note that cache invalidation is explicit + // via invalidate_range and need not be performed here. If using emulated + // mapping this must call iree_hal_buffer_emulated_unmap_range to release the + // transient resources. + (void)buffer; + iree_status_t status = iree_make_status(IREE_STATUS_UNIMPLEMENTED, + "buffer mapping not implemented"); + + return status; +} + +static iree_status_t iree_hal_null_buffer_invalidate_range( + iree_hal_buffer_t* base_buffer, iree_device_size_t local_byte_offset, + iree_device_size_t local_byte_length) { + iree_hal_null_buffer_t* buffer = iree_hal_null_buffer_cast(base_buffer); + + // TODO(null): invalidate the range if required by the buffer. Writes on the + // device are expected to be visible to the host after this returns. + (void)buffer; + iree_status_t status = iree_make_status(IREE_STATUS_UNIMPLEMENTED, + "buffer mapping not implemented"); + + return status; +} + +static iree_status_t iree_hal_null_buffer_flush_range( + iree_hal_buffer_t* base_buffer, iree_device_size_t local_byte_offset, + iree_device_size_t local_byte_length) { + iree_hal_null_buffer_t* buffer = iree_hal_null_buffer_cast(base_buffer); + + // TODO(null): flush the range if required by the buffer. Writes on the + // host are expected to be visible to the device after this returns. + (void)buffer; + iree_status_t status = iree_make_status(IREE_STATUS_UNIMPLEMENTED, + "buffer mapping not implemented"); + + return status; +} + +static const iree_hal_buffer_vtable_t iree_hal_null_buffer_vtable = { + .recycle = iree_hal_buffer_recycle, + .destroy = iree_hal_null_buffer_destroy, + .map_range = iree_hal_null_buffer_map_range, + .unmap_range = iree_hal_null_buffer_unmap_range, + .invalidate_range = iree_hal_null_buffer_invalidate_range, + .flush_range = iree_hal_null_buffer_flush_range, +}; diff --git a/runtime/src/iree/hal/drivers/null/buffer.h b/runtime/src/iree/hal/drivers/null/buffer.h new file mode 100644 index 000000000000..7e492f4d49d7 --- /dev/null +++ b/runtime/src/iree/hal/drivers/null/buffer.h @@ -0,0 +1,22 @@ +// Copyright 2024 The IREE Authors +// +// Licensed under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#ifndef IREE_HAL_DRIVERS_NULL_BUFFER_H_ +#define IREE_HAL_DRIVERS_NULL_BUFFER_H_ + +#include "iree/base/api.h" +#include "iree/hal/api.h" + +// Wraps a {Null} allocation in an iree_hal_buffer_t. +iree_status_t iree_hal_null_buffer_wrap( + iree_hal_allocator_t* allocator, iree_hal_memory_type_t memory_type, + iree_hal_memory_access_t allowed_access, + iree_hal_buffer_usage_t allowed_usage, iree_device_size_t allocation_size, + iree_device_size_t byte_offset, iree_device_size_t byte_length, + iree_hal_buffer_release_callback_t release_callback, + iree_allocator_t host_allocator, iree_hal_buffer_t** out_buffer); + +#endif // IREE_HAL_DRIVERS_NULL_BUFFER_H_ diff --git a/runtime/src/iree/hal/drivers/null/channel.c b/runtime/src/iree/hal/drivers/null/channel.c new file mode 100644 index 000000000000..195c3d5786cd --- /dev/null +++ b/runtime/src/iree/hal/drivers/null/channel.c @@ -0,0 +1,127 @@ +// Copyright 2024 The IREE Authors +// +// Licensed under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#include "iree/hal/drivers/null/channel.h" + +typedef struct iree_hal_null_channel_t { + iree_hal_resource_t resource; + iree_allocator_t host_allocator; + + // Parent channel this was split from, if any. + // This is only used to keep the parent channel live for as long as there are + // any split channels live (including transitive splits). + iree_hal_channel_t* parent_channel; +} iree_hal_null_channel_t; + +static const iree_hal_channel_vtable_t iree_hal_null_channel_vtable; + +static iree_hal_null_channel_t* iree_hal_null_channel_cast( + iree_hal_channel_t* base_value) { + IREE_HAL_ASSERT_TYPE(base_value, &iree_hal_null_channel_vtable); + return (iree_hal_null_channel_t*)base_value; +} + +static const iree_hal_null_channel_t* iree_hal_null_channel_const_cast( + const iree_hal_channel_t* base_value) { + IREE_HAL_ASSERT_TYPE(base_value, &iree_hal_null_channel_vtable); + return (const iree_hal_null_channel_t*)base_value; +} + +iree_status_t iree_hal_null_channel_create(iree_hal_channel_params_t params, + iree_allocator_t host_allocator, + iree_hal_channel_t** out_channel) { + IREE_ASSERT_ARGUMENT(out_channel); + *out_channel = NULL; + IREE_TRACE_ZONE_BEGIN(z0); + + iree_hal_null_channel_t* channel = NULL; + IREE_RETURN_AND_END_ZONE_IF_ERROR( + z0, iree_allocator_malloc(host_allocator, sizeof(*channel), + (void**)&channel)); + iree_hal_resource_initialize(&iree_hal_null_channel_vtable, + &channel->resource); + channel->host_allocator = host_allocator; + + // TODO(null): implement channel setup using params. Note that the id is not + // retained and must be copied local if needed beyond this function call. + iree_status_t status = iree_make_status( + IREE_STATUS_UNIMPLEMENTED, "collective channels not implemented"); + + if (iree_status_is_ok(status)) { + *out_channel = (iree_hal_channel_t*)channel; + } else { + iree_hal_channel_release((iree_hal_channel_t*)channel); + } + IREE_TRACE_ZONE_END(z0); + return status; +} + +static void iree_hal_null_channel_destroy(iree_hal_channel_t* base_channel) { + iree_hal_null_channel_t* channel = iree_hal_null_channel_cast(base_channel); + IREE_TRACE_ZONE_BEGIN(z0); + + iree_allocator_t host_allocator = channel->host_allocator; + + // TODO(null): destroy any implementation resources. + + iree_hal_channel_release(channel->parent_channel); + iree_allocator_free(host_allocator, channel); + + IREE_TRACE_ZONE_END(z0); +} + +static iree_status_t iree_hal_null_channel_split( + iree_hal_channel_t* base_channel, int32_t color, int32_t key, + iree_hal_channel_flags_t flags, iree_hal_channel_t** out_split_channel) { + iree_hal_null_channel_t* channel = iree_hal_null_channel_cast(base_channel); + + // TODO(null): split the channel and get any native resources required. + iree_status_t status = iree_make_status(IREE_STATUS_UNIMPLEMENTED, + "channel splitting not implemented"); + + // Wrap the split channel resources in a new HAL channel. + iree_hal_null_channel_t* split_channel = NULL; + if (iree_status_is_ok(status)) { + status = + iree_allocator_malloc(channel->host_allocator, sizeof(*split_channel), + (void**)&split_channel); + } + if (iree_status_is_ok(status)) { + iree_hal_resource_initialize(&iree_hal_null_channel_vtable, + &split_channel->resource); + split_channel->host_allocator = channel->host_allocator; + split_channel->parent_channel = base_channel; + iree_hal_channel_retain(base_channel); + + // TODO(null): transfer ownership of the implementation resources. + } + + if (iree_status_is_ok(status)) { + *out_split_channel = (iree_hal_channel_t*)split_channel; + } else { + iree_hal_channel_release((iree_hal_channel_t*)split_channel); + } + return status; +} + +static void iree_hal_null_channel_query_rank_and_count( + const iree_hal_channel_t* base_channel, int32_t* out_rank, + int32_t* out_count) { + const iree_hal_null_channel_t* channel = + iree_hal_null_channel_const_cast(base_channel); + + // TODO(null): query the rank and count from the implementation or cache them + // locally to avoid overheads (this may be called frequently). + (void)channel; + *out_rank = 0; + *out_count = 0; +} + +static const iree_hal_channel_vtable_t iree_hal_null_channel_vtable = { + .destroy = iree_hal_null_channel_destroy, + .split = iree_hal_null_channel_split, + .query_rank_and_count = iree_hal_null_channel_query_rank_and_count, +}; diff --git a/runtime/src/iree/hal/drivers/null/channel.h b/runtime/src/iree/hal/drivers/null/channel.h new file mode 100644 index 000000000000..83c4ef1aef88 --- /dev/null +++ b/runtime/src/iree/hal/drivers/null/channel.h @@ -0,0 +1,18 @@ +// Copyright 2024 The IREE Authors +// +// Licensed under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#ifndef IREE_HAL_DRIVERS_NULL_CHANNEL_H_ +#define IREE_HAL_DRIVERS_NULL_CHANNEL_H_ + +#include "iree/base/api.h" +#include "iree/hal/api.h" + +// Creates a {Null} HAL collective channel using the given |params|. +iree_status_t iree_hal_null_channel_create(iree_hal_channel_params_t params, + iree_allocator_t host_allocator, + iree_hal_channel_t** out_channel); + +#endif // IREE_HAL_DRIVERS_NULL_CHANNEL_H_ diff --git a/runtime/src/iree/hal/drivers/null/command_buffer.c b/runtime/src/iree/hal/drivers/null/command_buffer.c new file mode 100644 index 000000000000..9d474d44cb96 --- /dev/null +++ b/runtime/src/iree/hal/drivers/null/command_buffer.c @@ -0,0 +1,367 @@ +// Copyright 2024 The IREE Authors +// +// Licensed under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#include "iree/hal/drivers/null/command_buffer.h" + +#include "iree/hal/drivers/null/buffer.h" +#include "iree/hal/drivers/null/channel.h" +#include "iree/hal/drivers/null/executable.h" + +typedef struct iree_hal_null_command_buffer_t { + iree_hal_command_buffer_t base; + iree_allocator_t host_allocator; +} iree_hal_null_command_buffer_t; + +static const iree_hal_command_buffer_vtable_t + iree_hal_null_command_buffer_vtable; + +static iree_hal_null_command_buffer_t* iree_hal_null_command_buffer_cast( + iree_hal_command_buffer_t* base_value) { + IREE_HAL_ASSERT_TYPE(base_value, &iree_hal_null_command_buffer_vtable); + return (iree_hal_null_command_buffer_t*)base_value; +} + +iree_status_t iree_hal_null_command_buffer_create( + iree_hal_allocator_t* device_allocator, iree_hal_command_buffer_mode_t mode, + iree_hal_command_category_t command_categories, + iree_hal_queue_affinity_t queue_affinity, iree_host_size_t binding_capacity, + iree_allocator_t host_allocator, + iree_hal_command_buffer_t** out_command_buffer) { + IREE_ASSERT_ARGUMENT(out_command_buffer); + *out_command_buffer = NULL; + IREE_TRACE_ZONE_BEGIN(z0); + + iree_hal_null_command_buffer_t* command_buffer = NULL; + IREE_RETURN_AND_END_ZONE_IF_ERROR( + z0, + iree_allocator_malloc(host_allocator, + sizeof(*command_buffer) + + iree_hal_command_buffer_validation_state_size( + mode, binding_capacity), + (void**)&command_buffer)); + iree_hal_command_buffer_initialize( + device_allocator, mode, command_categories, queue_affinity, + binding_capacity, (uint8_t*)command_buffer + sizeof(*command_buffer), + &iree_hal_null_command_buffer_vtable, &command_buffer->base); + command_buffer->host_allocator = host_allocator; + + // TODO(null): allocate any additional resources for managing command buffer + // state. Some implementations may have their own command buffer/command list + // APIs this can route to or may need to implement it all themselves using + // iree_arena_t/block pools. Implementations should also retain any resources + // used during the recording and can use iree_hal_resource_set_t* to make that + // easier. + iree_status_t status = iree_make_status( + IREE_STATUS_UNIMPLEMENTED, "command buffers not yet implemented"); + + if (iree_status_is_ok(status)) { + *out_command_buffer = &command_buffer->base; + } else { + iree_hal_command_buffer_release(&command_buffer->base); + } + IREE_TRACE_ZONE_END(z0); + return status; +} + +static void iree_hal_null_command_buffer_destroy( + iree_hal_command_buffer_t* base_command_buffer) { + iree_hal_null_command_buffer_t* command_buffer = + iree_hal_null_command_buffer_cast(base_command_buffer); + iree_allocator_t host_allocator = command_buffer->host_allocator; + IREE_TRACE_ZONE_BEGIN(z0); + + // TODO(null): release any implementation resources and + // iree_hal_resource_set_t. + + iree_allocator_free(host_allocator, command_buffer); + + IREE_TRACE_ZONE_END(z0); +} + +bool iree_hal_null_command_buffer_isa( + iree_hal_command_buffer_t* command_buffer) { + return iree_hal_resource_is(&command_buffer->resource, + &iree_hal_null_command_buffer_vtable); +} + +static iree_status_t iree_hal_null_command_buffer_begin( + iree_hal_command_buffer_t* base_command_buffer) { + iree_hal_null_command_buffer_t* command_buffer = + iree_hal_null_command_buffer_cast(base_command_buffer); + + // TODO(null): if the implementation needs to route the begin to the + // implementation it can be done here. Note that creation may happen much + // earlier than recording and any expensive work should be deferred until this + // point to make profiling easier. + (void)command_buffer; + iree_status_t status = + iree_make_status(IREE_STATUS_UNIMPLEMENTED, + "command buffer recording start not implemented"); + + return status; +} + +static iree_status_t iree_hal_null_command_buffer_end( + iree_hal_command_buffer_t* base_command_buffer) { + iree_hal_null_command_buffer_t* command_buffer = + iree_hal_null_command_buffer_cast(base_command_buffer); + + // TODO(null): if recording requires multiple passes any fixup/linking can + // happen here. Recording-only resources are no longer needed after this point + // and can be disposed. + (void)command_buffer; + iree_status_t status = iree_make_status( + IREE_STATUS_UNIMPLEMENTED, "command buffer finalization not implemented"); + + return status; +} + +static void iree_hal_null_command_buffer_begin_debug_group( + iree_hal_command_buffer_t* base_command_buffer, iree_string_view_t label, + iree_hal_label_color_t label_color, + const iree_hal_label_location_t* location) { + iree_hal_null_command_buffer_t* command_buffer = + iree_hal_null_command_buffer_cast(base_command_buffer); + + // TODO(null): begin a nested debug group (push) if the implementation has a + // way to insert markers. This is informational and can be ignored. + (void)command_buffer; +} + +static void iree_hal_null_command_buffer_end_debug_group( + iree_hal_command_buffer_t* base_command_buffer) { + iree_hal_null_command_buffer_t* command_buffer = + iree_hal_null_command_buffer_cast(base_command_buffer); + + // TODO(null): end a nested debug group (pop). Always called 1:1 in stack + // order with begin_debug_group. + (void)command_buffer; +} + +static iree_status_t iree_hal_null_command_buffer_execution_barrier( + iree_hal_command_buffer_t* base_command_buffer, + iree_hal_execution_stage_t source_stage_mask, + iree_hal_execution_stage_t target_stage_mask, + iree_hal_execution_barrier_flags_t flags, + iree_host_size_t memory_barrier_count, + const iree_hal_memory_barrier_t* memory_barriers, + iree_host_size_t buffer_barrier_count, + const iree_hal_buffer_barrier_t* buffer_barriers) { + iree_hal_null_command_buffer_t* command_buffer = + iree_hal_null_command_buffer_cast(base_command_buffer); + + // TODO(null): barriers split the execution sequence into all operations that + // did happen before the barrier and all that will happen after. In + // implementations that have no concurrency this can be a no-op. This is + // effectively just a signal_event followed by a wait_event. + (void)command_buffer; + iree_status_t status = iree_make_status(IREE_STATUS_UNIMPLEMENTED, + "execution barriers not implemented"); + + return status; +} + +static iree_status_t iree_hal_null_command_buffer_signal_event( + iree_hal_command_buffer_t* base_command_buffer, iree_hal_event_t* event, + iree_hal_execution_stage_t source_stage_mask) { + iree_hal_null_command_buffer_t* command_buffer = + iree_hal_null_command_buffer_cast(base_command_buffer); + + // TODO(null): WIP API and may change; signals the given event allowing + // waiters to proceed. + (void)command_buffer; + iree_status_t status = + iree_make_status(IREE_STATUS_UNIMPLEMENTED, "events not implemented"); + + return status; +} + +static iree_status_t iree_hal_null_command_buffer_reset_event( + iree_hal_command_buffer_t* base_command_buffer, iree_hal_event_t* event, + iree_hal_execution_stage_t source_stage_mask) { + iree_hal_null_command_buffer_t* command_buffer = + iree_hal_null_command_buffer_cast(base_command_buffer); + + // TODO(null): WIP API and may change; resets the given event to unsignaled. + (void)command_buffer; + iree_status_t status = + iree_make_status(IREE_STATUS_UNIMPLEMENTED, "events not implemented"); + + return status; +} + +static iree_status_t iree_hal_null_command_buffer_wait_events( + iree_hal_command_buffer_t* base_command_buffer, + iree_host_size_t event_count, const iree_hal_event_t** events, + iree_hal_execution_stage_t source_stage_mask, + iree_hal_execution_stage_t target_stage_mask, + iree_host_size_t memory_barrier_count, + const iree_hal_memory_barrier_t* memory_barriers, + iree_host_size_t buffer_barrier_count, + const iree_hal_buffer_barrier_t* buffer_barriers) { + iree_hal_null_command_buffer_t* command_buffer = + iree_hal_null_command_buffer_cast(base_command_buffer); + + // TODO(null): WIP API and may change; waits on the list of events and enacts + // the specified set of barriers. Implementations without fine-grained + // tracking can treat this as an execution_barrier and ignore the + // memory/buffer barriers provided. + (void)command_buffer; + iree_status_t status = + iree_make_status(IREE_STATUS_UNIMPLEMENTED, "events not implemented"); + + return status; +} + +static iree_status_t iree_hal_null_command_buffer_discard_buffer( + iree_hal_command_buffer_t* base_command_buffer, + iree_hal_buffer_ref_t buffer_ref) { + iree_hal_null_command_buffer_t* command_buffer = + iree_hal_null_command_buffer_cast(base_command_buffer); + + // TODO(null): WIP API and may change; this is likely to become an + // madvise-like command that can be used to control prefetching and other + // cache behavior. The current discard behavior is a hint that the buffer + // contents will never be used again and that if they are in a cache they need + // not be written back to global memory. + (void)command_buffer; + iree_status_t status = iree_make_status(IREE_STATUS_UNIMPLEMENTED, + "discard buffer not implemented"); + + return status; +} + +static iree_status_t iree_hal_null_command_buffer_fill_buffer( + iree_hal_command_buffer_t* base_command_buffer, + iree_hal_buffer_ref_t target_ref, const void* pattern, + iree_host_size_t pattern_length) { + iree_hal_null_command_buffer_t* command_buffer = + iree_hal_null_command_buffer_cast(base_command_buffer); + + // TODO(null): memset on the buffer. The pattern_length is 1, 2, or 4 bytes. + // Note that the buffer may be a reference to a binding table slot in which + // case it will be provided during submission to a queue. + (void)command_buffer; + iree_status_t status = iree_make_status(IREE_STATUS_UNIMPLEMENTED, + "fill buffer not implemented"); + + return status; +} + +static iree_status_t iree_hal_null_command_buffer_update_buffer( + iree_hal_command_buffer_t* base_command_buffer, const void* source_buffer, + iree_host_size_t source_offset, iree_hal_buffer_ref_t target_ref) { + iree_hal_null_command_buffer_t* command_buffer = + iree_hal_null_command_buffer_cast(base_command_buffer); + + // TODO(null): embed and copy a small (~64KB) chunk of host memory to the + // target buffer. The source_buffer contents must be captured as they may + // change/be freed after this call completes. + // Note that the target buffer may be a reference to a binding table slot in + // which case it will be provided during submission to a queue. + (void)command_buffer; + iree_status_t status = iree_make_status(IREE_STATUS_UNIMPLEMENTED, + "update buffer not implemented"); + + return status; +} + +static iree_status_t iree_hal_null_command_buffer_copy_buffer( + iree_hal_command_buffer_t* base_command_buffer, + iree_hal_buffer_ref_t source_ref, iree_hal_buffer_ref_t target_ref) { + iree_hal_null_command_buffer_t* command_buffer = + iree_hal_null_command_buffer_cast(base_command_buffer); + + // TODO(null): memcpy between two buffers. The buffers must both be + // device-visible but may reside on either the host or device. + // Note that either buffer may be a reference to a binding table slot in + // which case it will be provided during submission to a queue. + (void)command_buffer; + iree_status_t status = iree_make_status(IREE_STATUS_UNIMPLEMENTED, + "copy buffer not implemented"); + + return status; +} + +static iree_status_t iree_hal_null_command_buffer_collective( + iree_hal_command_buffer_t* base_command_buffer, iree_hal_channel_t* channel, + iree_hal_collective_op_t op, uint32_t param, iree_hal_buffer_ref_t send_ref, + iree_hal_buffer_ref_t recv_ref, iree_device_size_t element_count) { + iree_hal_null_command_buffer_t* command_buffer = + iree_hal_null_command_buffer_cast(base_command_buffer); + + // TODO(null): perform the collective operation defined by op. See the headers + // for more information. The channel is fixed for a particular recording but + // note that either buffer may be a reference to a binding table slot in + // which case it will be provided during submission to a queue. + (void)command_buffer; + iree_status_t status = iree_make_status(IREE_STATUS_UNIMPLEMENTED, + "collectives not implemented"); + + return status; +} + +static iree_status_t iree_hal_null_command_buffer_dispatch( + iree_hal_command_buffer_t* base_command_buffer, + iree_hal_executable_t* executable, int32_t entry_point, + const uint32_t workgroup_count[3], iree_const_byte_span_t constants, + iree_hal_buffer_ref_list_t bindings, iree_hal_dispatch_flags_t flags) { + iree_hal_null_command_buffer_t* command_buffer = + iree_hal_null_command_buffer_cast(base_command_buffer); + + // TODO(null): dispatch the specified executable entry point with the given + // workgroup count. The constants must be copied into the command buffer as + // they may be mutated or freed after this call returns. + // Note that any of the bindings may be references to binding table slots in + // which case they will be provided during submission to a queue. + (void)command_buffer; + iree_status_t status = + iree_make_status(IREE_STATUS_UNIMPLEMENTED, "dispatch not implemented"); + + return status; +} + +static iree_status_t iree_hal_null_command_buffer_dispatch_indirect( + iree_hal_command_buffer_t* base_command_buffer, + iree_hal_executable_t* executable, int32_t entry_point, + iree_hal_buffer_ref_t workgroups_ref, iree_const_byte_span_t constants, + iree_hal_buffer_ref_list_t bindings, iree_hal_dispatch_flags_t flags) { + iree_hal_null_command_buffer_t* command_buffer = + iree_hal_null_command_buffer_cast(base_command_buffer); + + // TODO(null): dispatch the specified executable entry point with a workgroup + // count that is stored in the given workgroup count buffer as a uint32_t[3]. + // The workgroup count may change up until immediately prior to the dispatch. + // The constants must be copied into the command buffer as they may be mutated + // or freed after this call returns. Note that any of the bindings may be + // references to binding table slots in which case they will be provided + // during submission to a queue. + (void)command_buffer; + iree_status_t status = iree_make_status(IREE_STATUS_UNIMPLEMENTED, + "indirect dispatch not implemented"); + + return status; +} + +static const iree_hal_command_buffer_vtable_t + iree_hal_null_command_buffer_vtable = { + .destroy = iree_hal_null_command_buffer_destroy, + .begin = iree_hal_null_command_buffer_begin, + .end = iree_hal_null_command_buffer_end, + .begin_debug_group = iree_hal_null_command_buffer_begin_debug_group, + .end_debug_group = iree_hal_null_command_buffer_end_debug_group, + .execution_barrier = iree_hal_null_command_buffer_execution_barrier, + .signal_event = iree_hal_null_command_buffer_signal_event, + .reset_event = iree_hal_null_command_buffer_reset_event, + .wait_events = iree_hal_null_command_buffer_wait_events, + .discard_buffer = iree_hal_null_command_buffer_discard_buffer, + .fill_buffer = iree_hal_null_command_buffer_fill_buffer, + .update_buffer = iree_hal_null_command_buffer_update_buffer, + .copy_buffer = iree_hal_null_command_buffer_copy_buffer, + .collective = iree_hal_null_command_buffer_collective, + .dispatch = iree_hal_null_command_buffer_dispatch, + .dispatch_indirect = iree_hal_null_command_buffer_dispatch_indirect, +}; diff --git a/runtime/src/iree/hal/drivers/null/command_buffer.h b/runtime/src/iree/hal/drivers/null/command_buffer.h new file mode 100644 index 000000000000..cca92367dd82 --- /dev/null +++ b/runtime/src/iree/hal/drivers/null/command_buffer.h @@ -0,0 +1,25 @@ +// Copyright 2024 The IREE Authors +// +// Licensed under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#ifndef IREE_HAL_DRIVERS_NULL_COMMAND_BUFFER_H_ +#define IREE_HAL_DRIVERS_NULL_COMMAND_BUFFER_H_ + +#include "iree/base/api.h" +#include "iree/hal/api.h" + +// Creates {Null} command buffer. +iree_status_t iree_hal_null_command_buffer_create( + iree_hal_allocator_t* device_allocator, iree_hal_command_buffer_mode_t mode, + iree_hal_command_category_t command_categories, + iree_hal_queue_affinity_t queue_affinity, iree_host_size_t binding_capacity, + iree_allocator_t host_allocator, + iree_hal_command_buffer_t** out_command_buffer); + +// Returns true if |command_buffer| is a {Null} command buffer. +bool iree_hal_null_command_buffer_isa( + iree_hal_command_buffer_t* command_buffer); + +#endif // IREE_HAL_DRIVERS_NULL_COMMAND_BUFFER_H_ diff --git a/runtime/src/iree/hal/drivers/null/device.c b/runtime/src/iree/hal/drivers/null/device.c new file mode 100644 index 000000000000..aaa7b1591fd5 --- /dev/null +++ b/runtime/src/iree/hal/drivers/null/device.c @@ -0,0 +1,559 @@ +// Copyright 2024 The IREE Authors +// +// Licensed under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#include "iree/hal/drivers/null/device.h" + +#include "iree/hal/drivers/null/allocator.h" +#include "iree/hal/drivers/null/api.h" +#include "iree/hal/drivers/null/channel.h" +#include "iree/hal/drivers/null/command_buffer.h" +#include "iree/hal/drivers/null/event.h" +#include "iree/hal/drivers/null/executable.h" +#include "iree/hal/drivers/null/executable_cache.h" +#include "iree/hal/drivers/null/semaphore.h" +#include "iree/hal/utils/file_transfer.h" +#include "iree/hal/utils/memory_file.h" + +typedef struct iree_hal_null_device_t { + iree_hal_resource_t resource; + iree_string_view_t identifier; + + iree_allocator_t host_allocator; + iree_hal_allocator_t* device_allocator; + + // Optional provider used for creating/configuring collective channels. + iree_hal_channel_provider_t* channel_provider; + + // + trailing identifier string storage +} iree_hal_null_device_t; + +static const iree_hal_device_vtable_t iree_hal_null_device_vtable; + +static iree_hal_null_device_t* iree_hal_null_device_cast( + iree_hal_device_t* base_value) { + IREE_HAL_ASSERT_TYPE(base_value, &iree_hal_null_device_vtable); + return (iree_hal_null_device_t*)base_value; +} + +void iree_hal_null_device_options_initialize( + iree_hal_null_device_options_t* out_options) { + memset(out_options, 0, sizeof(*out_options)); + // TODO(null): set defaults based on compiler configuration. Flags should not + // be used as multiple devices may be configured within the process or the + // hosting application may be authored in python/etc that does not use a flags + // mechanism accessible here. +} + +static iree_status_t iree_hal_null_device_options_verify( + const iree_hal_null_device_options_t* options) { + // TODO(null): verify that the parameters are within expected ranges and any + // requested features are supported. + return iree_ok_status(); +} + +iree_status_t iree_hal_null_device_create( + iree_string_view_t identifier, + const iree_hal_null_device_options_t* options, + iree_allocator_t host_allocator, iree_hal_device_t** out_device) { + IREE_ASSERT_ARGUMENT(options); + IREE_ASSERT_ARGUMENT(out_device); + *out_device = NULL; + IREE_TRACE_ZONE_BEGIN(z0); + + // Verify the parameters prior to creating resources. + IREE_RETURN_AND_END_ZONE_IF_ERROR( + z0, iree_hal_null_device_options_verify(options)); + + iree_hal_null_device_t* device = NULL; + iree_host_size_t total_size = sizeof(*device) + identifier.size; + IREE_RETURN_AND_END_ZONE_IF_ERROR( + z0, iree_allocator_malloc(host_allocator, total_size, (void**)&device)); + iree_hal_resource_initialize(&iree_hal_null_device_vtable, &device->resource); + iree_string_view_append_to_buffer( + identifier, &device->identifier, + (char*)device + total_size - identifier.size); + device->host_allocator = host_allocator; + + // TODO(null): pass device handles and pool configuration to the allocator. + // Some implementations may share allocators across multiple devices created + // from the same driver. + iree_status_t status = + iree_hal_null_allocator_create(host_allocator, &device->device_allocator); + + if (iree_status_is_ok(status)) { + *out_device = (iree_hal_device_t*)device; + } else { + iree_hal_device_release((iree_hal_device_t*)device); + } + IREE_TRACE_ZONE_END(z0); + return status; +} + +static void iree_hal_null_device_destroy(iree_hal_device_t* base_device) { + iree_hal_null_device_t* device = iree_hal_null_device_cast(base_device); + iree_allocator_t host_allocator = iree_hal_device_host_allocator(base_device); + IREE_TRACE_ZONE_BEGIN(z0); + + // TODO(null): release all implementation resources here. It's expected that + // this is only called once all outstanding resources created with this device + // have been released by the application and no work is outstanding. If the + // implementation performs internal async operations those should be shutdown + // and joined first. + + iree_hal_allocator_release(device->device_allocator); + iree_hal_channel_provider_release(device->channel_provider); + + iree_allocator_free(host_allocator, device); + + IREE_TRACE_ZONE_END(z0); +} + +static iree_string_view_t iree_hal_null_device_id( + iree_hal_device_t* base_device) { + iree_hal_null_device_t* device = iree_hal_null_device_cast(base_device); + return device->identifier; +} + +static iree_allocator_t iree_hal_null_device_host_allocator( + iree_hal_device_t* base_device) { + iree_hal_null_device_t* device = iree_hal_null_device_cast(base_device); + return device->host_allocator; +} + +static iree_hal_allocator_t* iree_hal_null_device_allocator( + iree_hal_device_t* base_device) { + iree_hal_null_device_t* device = iree_hal_null_device_cast(base_device); + return device->device_allocator; +} + +static void iree_hal_null_replace_device_allocator( + iree_hal_device_t* base_device, iree_hal_allocator_t* new_allocator) { + iree_hal_null_device_t* device = iree_hal_null_device_cast(base_device); + iree_hal_allocator_retain(new_allocator); + iree_hal_allocator_release(device->device_allocator); + device->device_allocator = new_allocator; +} + +static void iree_hal_null_replace_channel_provider( + iree_hal_device_t* base_device, iree_hal_channel_provider_t* new_provider) { + iree_hal_null_device_t* device = iree_hal_null_device_cast(base_device); + iree_hal_channel_provider_retain(new_provider); + iree_hal_channel_provider_release(device->channel_provider); + device->channel_provider = new_provider; +} + +static iree_status_t iree_hal_null_device_trim(iree_hal_device_t* base_device) { + iree_hal_null_device_t* device = iree_hal_null_device_cast(base_device); + + // TODO(null): if the device has any cached resources that can be safely + // dropped here (unused pools/etc). This is usually called in low-memory + // situations or when the HAL device will not be used for awhile (device + // entering sleep mode or a low power state, etc). + + IREE_RETURN_IF_ERROR(iree_hal_allocator_trim(device->device_allocator)); + + return iree_ok_status(); +} + +static iree_status_t iree_hal_null_device_query_i64( + iree_hal_device_t* base_device, iree_string_view_t category, + iree_string_view_t key, int64_t* out_value) { + iree_hal_null_device_t* device = iree_hal_null_device_cast(base_device); + *out_value = 0; + + // TODO(null): implement additional queries. These are stubs for common ones + // as used by the compiler. Targets may have their own, though, and connect + // with them by emitting `hal.device.query` ops in programs or calling the + // query method at runtime via the HAL API. + + if (iree_string_view_equal(category, IREE_SV("hal.device.id"))) { + // NOTE: this is a fuzzy match and can allow a program to work with multiple + // device implementations. + *out_value = + iree_string_view_match_pattern(device->identifier, key) ? 1 : 0; + return iree_ok_status(); + } + + if (iree_string_view_equal(category, IREE_SV("hal.executable.format"))) { + // NOTE: this is a fuzzy match and can allow multiple formats to be used + // (this should return 1 for any format supported). + // TODO(null): match a format and return true. + *out_value = 0; + return iree_ok_status(); + } + + // TODO(null): return basic queries for concurrency to allow programs to + // estimate potential utilization. + if (iree_string_view_equal(category, IREE_SV("hal.device"))) { + if (iree_string_view_equal(key, IREE_SV("concurrency"))) { + *out_value = 1; + return iree_ok_status(); + } + } else if (iree_string_view_equal(category, IREE_SV("hal.dispatch"))) { + if (iree_string_view_equal(key, IREE_SV("concurrency"))) { + *out_value = 1; + return iree_ok_status(); + } + } + + return iree_make_status( + IREE_STATUS_NOT_FOUND, + "unknown device configuration key value '%.*s :: %.*s'", + (int)category.size, category.data, (int)key.size, key.data); +} + +static iree_status_t iree_hal_null_device_create_channel( + iree_hal_device_t* base_device, iree_hal_queue_affinity_t queue_affinity, + iree_hal_channel_params_t params, iree_hal_channel_t** out_channel) { + iree_hal_null_device_t* device = iree_hal_null_device_cast(base_device); + + // TODO(null): pass any additional resources required to create the channel. + // The device->channel_provider can be used to get default rank/count, + // exchange IDs, etc as needed. + (void)device; + + return iree_hal_null_channel_create( + params, iree_hal_device_host_allocator(base_device), out_channel); +} + +static iree_status_t iree_hal_null_device_create_command_buffer( + iree_hal_device_t* base_device, iree_hal_command_buffer_mode_t mode, + iree_hal_command_category_t command_categories, + iree_hal_queue_affinity_t queue_affinity, iree_host_size_t binding_capacity, + iree_hal_command_buffer_t** out_command_buffer) { + iree_hal_null_device_t* device = iree_hal_null_device_cast(base_device); + // TODO(null): pass any additional resources required to create the command + // buffer. The implementation could pool command buffers here. + return iree_hal_null_command_buffer_create( + iree_hal_device_allocator(base_device), mode, command_categories, + queue_affinity, binding_capacity, device->host_allocator, + out_command_buffer); +} + +static iree_status_t iree_hal_null_device_create_event( + iree_hal_device_t* base_device, iree_hal_queue_affinity_t queue_affinity, + iree_hal_event_flags_t flags, iree_hal_event_t** out_event) { + iree_hal_null_device_t* device = iree_hal_null_device_cast(base_device); + + // TODO(null): pass any additional resources required to create the event. + // The implementation could pool events here. + (void)device; + + return iree_hal_null_event_create(queue_affinity, flags, + iree_hal_device_host_allocator(base_device), + out_event); +} + +static iree_status_t iree_hal_null_device_create_executable_cache( + iree_hal_device_t* base_device, iree_string_view_t identifier, + iree_loop_t loop, iree_hal_executable_cache_t** out_executable_cache) { + iree_hal_null_device_t* device = iree_hal_null_device_cast(base_device); + + // TODO(null): pass any additional resources required during executable + // creation or cache management. + (void)device; + + return iree_hal_null_executable_cache_create( + identifier, iree_hal_device_host_allocator(base_device), + out_executable_cache); +} + +static iree_status_t iree_hal_null_device_import_file( + iree_hal_device_t* base_device, iree_hal_queue_affinity_t queue_affinity, + iree_hal_memory_access_t access, iree_io_file_handle_t* handle, + iree_hal_external_file_flags_t flags, iree_hal_file_t** out_file) { + // TODO(null): if the implementation supports native file operations + // definitely prefer that. The emulated file I/O present here as a default is + // inefficient. The queue affinity specifies which queues may access the file + // via read and write queue operations. + if (iree_io_file_handle_type(handle) != + IREE_IO_FILE_HANDLE_TYPE_HOST_ALLOCATION) { + return iree_make_status( + IREE_STATUS_UNAVAILABLE, + "implementation does not support the external file type"); + } + return iree_hal_memory_file_wrap( + queue_affinity, access, handle, iree_hal_device_allocator(base_device), + iree_hal_device_host_allocator(base_device), out_file); +} + +static iree_status_t iree_hal_null_device_create_semaphore( + iree_hal_device_t* base_device, uint64_t initial_value, + iree_hal_semaphore_flags_t flags, iree_hal_semaphore_t** out_semaphore) { + iree_hal_null_device_t* device = iree_hal_null_device_cast(base_device); + + // TODO(null): pass any additional resources required to create or track the + // semaphore. The implementation could pool semaphores here. + (void)device; + + return iree_hal_null_semaphore_create(initial_value, flags, + device->host_allocator, out_semaphore); +} + +static iree_hal_semaphore_compatibility_t +iree_hal_null_device_query_semaphore_compatibility( + iree_hal_device_t* base_device, iree_hal_semaphore_t* semaphore) { + iree_hal_null_device_t* device = iree_hal_null_device_cast(base_device); + + // TODO(null): return the appropriate bits for the provided semaphore + // indicating how it may be used with this device. The semaphore may have been + // created or imported on this device or any other device from the same + // driver. Certain implementations may allow semaphores from other drivers to + // be used and those can be checked here (though the API to do this isn't + // implemented yet). + (void)device; + iree_hal_semaphore_compatibility_t compatibility = + IREE_HAL_SEMAPHORE_COMPATIBILITY_NONE; + + return compatibility; +} + +static iree_status_t iree_hal_null_device_queue_alloca( + iree_hal_device_t* base_device, iree_hal_queue_affinity_t queue_affinity, + const iree_hal_semaphore_list_t wait_semaphore_list, + const iree_hal_semaphore_list_t signal_semaphore_list, + iree_hal_allocator_pool_t pool, iree_hal_buffer_params_t params, + iree_device_size_t allocation_size, + iree_hal_buffer_t** IREE_RESTRICT out_buffer) { + iree_hal_null_device_t* device = iree_hal_null_device_cast(base_device); + + // TODO(null): perform an allocation of a transient buffer in queue order. + // The allocation may be used on any queue set in the provided queue affinity. + // Deallocation via queue_dealloca is preferred but users are allowed to + // deallocate the buffer on the host via iree_hal_buffer_release even if they + // allocated it with queue_alloca. + (void)device; + iree_status_t status = iree_make_status(IREE_STATUS_UNIMPLEMENTED, + "queue alloca not implemented"); + + return status; +} + +static iree_status_t iree_hal_null_device_queue_dealloca( + iree_hal_device_t* base_device, iree_hal_queue_affinity_t queue_affinity, + const iree_hal_semaphore_list_t wait_semaphore_list, + const iree_hal_semaphore_list_t signal_semaphore_list, + iree_hal_buffer_t* buffer) { + iree_hal_null_device_t* device = iree_hal_null_device_cast(base_device); + + // TODO(null): perform a deallocation of the transient buffer in queue order. + // Only buffers allocated with queue_alloca on the same device will be passed. + // Note that different queues on the same device may have allocated the buffer + // and if the same queue must deallocate it the implementation will need to + // track that on the buffer. The user is allowed to deallocate the buffer on + // the host via iree_hal_buffer_release even if they allocated it with + // queue_alloca. + (void)device; + iree_status_t status = iree_make_status(IREE_STATUS_UNIMPLEMENTED, + "queue dealloca not implemented"); + + return status; +} + +static iree_status_t iree_hal_null_device_queue_read( + iree_hal_device_t* base_device, iree_hal_queue_affinity_t queue_affinity, + const iree_hal_semaphore_list_t wait_semaphore_list, + const iree_hal_semaphore_list_t signal_semaphore_list, + iree_hal_file_t* source_file, uint64_t source_offset, + iree_hal_buffer_t* target_buffer, iree_device_size_t target_offset, + iree_device_size_t length, uint32_t flags) { + // TODO(null): if native support for file operations are available then + // definitely prefer those over the emulated implementation provided here by + // default. The implementation performs allocations, creates semaphores, and + // submits command buffers with host-device blocking behavior. + + // TODO: expose streaming chunk count/size options. + iree_status_t loop_status = iree_ok_status(); + iree_hal_file_transfer_options_t options = { + .loop = iree_loop_inline(&loop_status), + .chunk_count = IREE_HAL_FILE_TRANSFER_CHUNK_COUNT_DEFAULT, + .chunk_size = IREE_HAL_FILE_TRANSFER_CHUNK_SIZE_DEFAULT, + }; + IREE_RETURN_IF_ERROR(iree_hal_device_queue_read_streaming( + base_device, queue_affinity, wait_semaphore_list, signal_semaphore_list, + source_file, source_offset, target_buffer, target_offset, length, flags, + options)); + return loop_status; +} + +static iree_status_t iree_hal_null_device_queue_write( + iree_hal_device_t* base_device, iree_hal_queue_affinity_t queue_affinity, + const iree_hal_semaphore_list_t wait_semaphore_list, + const iree_hal_semaphore_list_t signal_semaphore_list, + iree_hal_buffer_t* source_buffer, iree_device_size_t source_offset, + iree_hal_file_t* target_file, uint64_t target_offset, + iree_device_size_t length, uint32_t flags) { + // TODO(null): if native support for file operations are available then + // definitely prefer those over the emulated implementation provided here by + // default. The implementation performs allocations, creates semaphores, and + // submits command buffers with host-device blocking behavior. + + // TODO: expose streaming chunk count/size options. + iree_status_t loop_status = iree_ok_status(); + iree_hal_file_transfer_options_t options = { + .loop = iree_loop_inline(&loop_status), + .chunk_count = IREE_HAL_FILE_TRANSFER_CHUNK_COUNT_DEFAULT, + .chunk_size = IREE_HAL_FILE_TRANSFER_CHUNK_SIZE_DEFAULT, + }; + IREE_RETURN_IF_ERROR(iree_hal_device_queue_write_streaming( + base_device, queue_affinity, wait_semaphore_list, signal_semaphore_list, + source_buffer, source_offset, target_file, target_offset, length, flags, + options)); + return loop_status; +} + +static iree_status_t iree_hal_null_device_queue_execute( + iree_hal_device_t* base_device, iree_hal_queue_affinity_t queue_affinity, + const iree_hal_semaphore_list_t wait_semaphore_list, + const iree_hal_semaphore_list_t signal_semaphore_list, + iree_host_size_t command_buffer_count, + iree_hal_command_buffer_t* const* command_buffers, + iree_hal_buffer_binding_table_t const* binding_tables) { + iree_hal_null_device_t* device = iree_hal_null_device_cast(base_device); + + // TODO(null): implement a wait, execute, and signal queue operation. The + // queue affinity can be used to determine which top-level execution resources + // are to be used when executing and it can be assumed that all resources + // required for execution are accessible on those queues. If more than one + // queue is specified the implementation may use any it prefers from the set. + // If more than one command buffer is provided it is expected that they are + // executed in order on the same queue. + + // TODO(null): optional binding tables matching 1:1 with the command buffers + // are provided for any indirect command buffers (those who have a + // binding_capacity > 0). The binding tables must be captured by the + // implementation as they may be mutated or freed by the caller immediately + // after this call returns. + + // TODO(null): do this async - callers may be submitting work to multiple + // devices or queues on the same device from the same thread and blocking here + // will prevent both concurrency and pipelining. + + (void)device; + iree_status_t status = iree_make_status(IREE_STATUS_UNIMPLEMENTED, + "queue execute not implemented"); + + return status; +} + +static iree_status_t iree_hal_null_device_queue_flush( + iree_hal_device_t* base_device, iree_hal_queue_affinity_t queue_affinity) { + iree_hal_null_device_t* device = iree_hal_null_device_cast(base_device); + + // TODO(null): though buffering queue operations is not recommended if any + // buffering has been performed it must be flushed here. Callers may be + // indicating that they are about to suspend themselves waiting for submitted + // work to complete. + (void)device; + iree_status_t status = iree_make_status(IREE_STATUS_UNIMPLEMENTED, + "queue flush not implemented"); + + return status; +} + +static iree_status_t iree_hal_null_device_wait_semaphores( + iree_hal_device_t* base_device, iree_hal_wait_mode_t wait_mode, + const iree_hal_semaphore_list_t semaphore_list, iree_timeout_t timeout) { + iree_hal_null_device_t* device = iree_hal_null_device_cast(base_device); + + // TODO(null): implement multi-wait as either an ALL (AND) or ANY (OR) + // operation. Semaphores are expected to be compatible with the device today + // and may come from other device instances provided by the same driver or + // have been imported by a device instance. + + // TODO(null): if any semaphore has a failure status set return + // `iree_status_from_code(IREE_STATUS_ABORTED)`. Avoid a full status as it may + // capture a backtrace and allocate and callers are expected to follow up a + // failed wait with a query to get the status. + + // TODO(null): prefer having a fast-path for if the semaphores are + // known-signaled in user-mode. This can usually avoid syscalls/ioctls and + // potential context switches in polling cases. + + // TODO(null): check for `iree_timeout_is_immediate(timeout)` and return + // immediately if the condition is not satisfied before waiting with + // `iree_status_from_code(IREE_STATUS_DEADLINE_EXCEEDED)`. Prefer the raw code + // status instead of a full status object as immediate timeouts are used when + // polling and a full status may capture a backtrace and allocate. + + (void)device; + iree_status_t status = iree_make_status( + IREE_STATUS_UNIMPLEMENTED, "semaphore multi-wait not implemented"); + + return status; +} + +static iree_status_t iree_hal_null_device_profiling_begin( + iree_hal_device_t* base_device, + const iree_hal_device_profiling_options_t* options) { + iree_hal_null_device_t* device = iree_hal_null_device_cast(base_device); + + // TODO(null): set implementation-defined device or global profiling modes. + // This will be paired with a profiling_end call at some point in the future. + // Hosting applications may periodically call profiling_flush. + (void)device; + iree_status_t status = iree_make_status(IREE_STATUS_UNIMPLEMENTED, + "device profiling not implemented"); + + return status; +} + +static iree_status_t iree_hal_null_device_profiling_flush( + iree_hal_device_t* base_device) { + iree_hal_null_device_t* device = iree_hal_null_device_cast(base_device); + + // TODO(null): flush if needed. May be no-op. Any accumulated profiling + // information should be carried across the flush but the event can be used to + // reclaim resources or perform other expensive bookkeeping. Benchmarks, for + // example, are expected to call this periodically with their timing + // suspended. + (void)device; + iree_status_t status = iree_make_status(IREE_STATUS_UNIMPLEMENTED, + "device profiling not implemented"); + + return status; +} + +static iree_status_t iree_hal_null_device_profiling_end( + iree_hal_device_t* base_device) { + iree_hal_null_device_t* device = iree_hal_null_device_cast(base_device); + + // TODO(null): unset whatever profiling_begin set, if anything. May be no-op. + (void)device; + iree_status_t status = iree_make_status(IREE_STATUS_UNIMPLEMENTED, + "device profiling not implemented"); + + return status; +} + +static const iree_hal_device_vtable_t iree_hal_null_device_vtable = { + .destroy = iree_hal_null_device_destroy, + .id = iree_hal_null_device_id, + .host_allocator = iree_hal_null_device_host_allocator, + .device_allocator = iree_hal_null_device_allocator, + .replace_device_allocator = iree_hal_null_replace_device_allocator, + .replace_channel_provider = iree_hal_null_replace_channel_provider, + .trim = iree_hal_null_device_trim, + .query_i64 = iree_hal_null_device_query_i64, + .create_channel = iree_hal_null_device_create_channel, + .create_command_buffer = iree_hal_null_device_create_command_buffer, + .create_event = iree_hal_null_device_create_event, + .create_executable_cache = iree_hal_null_device_create_executable_cache, + .import_file = iree_hal_null_device_import_file, + .create_semaphore = iree_hal_null_device_create_semaphore, + .query_semaphore_compatibility = + iree_hal_null_device_query_semaphore_compatibility, + .queue_alloca = iree_hal_null_device_queue_alloca, + .queue_dealloca = iree_hal_null_device_queue_dealloca, + .queue_read = iree_hal_null_device_queue_read, + .queue_write = iree_hal_null_device_queue_write, + .queue_execute = iree_hal_null_device_queue_execute, + .queue_flush = iree_hal_null_device_queue_flush, + .wait_semaphores = iree_hal_null_device_wait_semaphores, + .profiling_begin = iree_hal_null_device_profiling_begin, + .profiling_flush = iree_hal_null_device_profiling_flush, + .profiling_end = iree_hal_null_device_profiling_end, +}; diff --git a/runtime/src/iree/hal/drivers/null/device.h b/runtime/src/iree/hal/drivers/null/device.h new file mode 100644 index 000000000000..aa70db6408d6 --- /dev/null +++ b/runtime/src/iree/hal/drivers/null/device.h @@ -0,0 +1,17 @@ +// Copyright 2024 The IREE Authors +// +// Licensed under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#ifndef IREE_HAL_DRIVERS_NULL_DEVICE_H_ +#define IREE_HAL_DRIVERS_NULL_DEVICE_H_ + +#include "iree/base/api.h" +#include "iree/hal/api.h" + +// NOTE: nothing in the skeleton implementation. Device creation and adoption is +// part of the public API header. This header can contain internal types and +// functions. + +#endif // IREE_HAL_DRIVERS_NULL_DEVICE_H_ diff --git a/runtime/src/iree/hal/drivers/null/driver.c b/runtime/src/iree/hal/drivers/null/driver.c new file mode 100644 index 000000000000..94be18a45364 --- /dev/null +++ b/runtime/src/iree/hal/drivers/null/driver.c @@ -0,0 +1,201 @@ +// Copyright 2024 The IREE Authors +// +// Licensed under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#include "iree/hal/drivers/null/driver.h" + +#include "iree/hal/drivers/null/api.h" + +// TODO(null): if it's possible to have more than one device use real IDs. +// This is a placeholder for this skeleton that just indicates the first and +// only device. +#define IREE_HAL_NULL_DEVICE_ID_DEFAULT 0 + +typedef struct iree_hal_null_driver_t { + iree_hal_resource_t resource; + iree_allocator_t host_allocator; + + iree_string_view_t identifier; + iree_hal_null_driver_options_t options; + + // + trailing identifier string storage +} iree_hal_null_driver_t; + +static const iree_hal_driver_vtable_t iree_hal_null_driver_vtable; + +static iree_hal_null_driver_t* iree_hal_null_driver_cast( + iree_hal_driver_t* base_value) { + IREE_HAL_ASSERT_TYPE(base_value, &iree_hal_null_driver_vtable); + return (iree_hal_null_driver_t*)base_value; +} + +void iree_hal_null_driver_options_initialize( + iree_hal_null_driver_options_t* out_options) { + memset(out_options, 0, sizeof(*out_options)); + + // TODO(null): set defaults based on compiler configuration. Flags should not + // be used as multiple devices may be configured within the process or the + // hosting application may be authored in python/etc that does not use a flags + // mechanism accessible here. + + iree_hal_null_device_options_initialize(&out_options->default_device_options); +} + +static iree_status_t iree_hal_null_driver_options_verify( + const iree_hal_null_driver_options_t* options) { + // TODO(null): verify that the parameters are within expected ranges and any + // requested features are supported. + + return iree_ok_status(); +} + +IREE_API_EXPORT iree_status_t iree_hal_null_driver_create( + iree_string_view_t identifier, + const iree_hal_null_driver_options_t* options, + iree_allocator_t host_allocator, iree_hal_driver_t** out_driver) { + IREE_ASSERT_ARGUMENT(options); + IREE_ASSERT_ARGUMENT(out_driver); + *out_driver = NULL; + IREE_TRACE_ZONE_BEGIN(z0); + + // TODO(null): verify options; this may be moved after any libraries are + // loaded so the verification can use underlying implementation queries. + IREE_RETURN_AND_END_ZONE_IF_ERROR( + z0, iree_hal_null_driver_options_verify(options)); + + iree_hal_null_driver_t* driver = NULL; + iree_host_size_t total_size = sizeof(*driver) + identifier.size; + IREE_RETURN_AND_END_ZONE_IF_ERROR( + z0, iree_allocator_malloc(host_allocator, total_size, (void**)&driver)); + iree_hal_resource_initialize(&iree_hal_null_driver_vtable, &driver->resource); + driver->host_allocator = host_allocator; + iree_string_view_append_to_buffer( + identifier, &driver->identifier, + (char*)driver + total_size - identifier.size); + + // TODO(null): if there are any string fields then they will need to be + // retained as well (similar to the identifier they can be tagged on to the + // end of the driver struct). + memcpy(&driver->options, options, sizeof(*options)); + + // TODO(null): load libraries and query driver support from the system. + // Devices need not be enumerated here if doing so is expensive; the + // application may create drivers just to see if they are present but defer + // device enumeration until the user requests one. Underlying implementations + // can sometimes do bonkers static init stuff as soon as they are touched and + // this code may want to do that on-demand instead. + iree_status_t status = + iree_make_status(IREE_STATUS_UNIMPLEMENTED, "driver not implemented"); + + if (iree_status_is_ok(status)) { + *out_driver = (iree_hal_driver_t*)driver; + } else { + iree_hal_driver_release((iree_hal_driver_t*)driver); + } + IREE_TRACE_ZONE_END(z0); + return status; +} + +static void iree_hal_null_driver_destroy(iree_hal_driver_t* base_driver) { + iree_hal_null_driver_t* driver = iree_hal_null_driver_cast(base_driver); + iree_allocator_t host_allocator = driver->host_allocator; + IREE_TRACE_ZONE_BEGIN(z0); + + // TODO(null): if the driver loaded any libraries they should be closed here. + + iree_allocator_free(host_allocator, driver); + + IREE_TRACE_ZONE_END(z0); +} + +static iree_status_t iree_hal_null_driver_query_available_devices( + iree_hal_driver_t* base_driver, iree_allocator_t host_allocator, + iree_host_size_t* out_device_info_count, + iree_hal_device_info_t** out_device_infos) { + // TODO(null): query available devices and populate the output. Note that + // unlike most IREE functions this allocates if required in order to allow + // this to return uncached information. Uncached is preferred as it allows + // devices that may come and go (power toggles, user visibility toggles, etc) + // through a process lifetime to appear without needing a full restart. + static const iree_hal_device_info_t device_infos[1] = { + { + .device_id = IREE_HAL_NULL_DEVICE_ID_DEFAULT, + .name = iree_string_view_literal("default"), + }, + }; + *out_device_info_count = IREE_ARRAYSIZE(device_infos); + return iree_allocator_clone( + host_allocator, + iree_make_const_byte_span(device_infos, sizeof(device_infos)), + (void**)out_device_infos); +} + +static iree_status_t iree_hal_null_driver_dump_device_info( + iree_hal_driver_t* base_driver, iree_hal_device_id_t device_id, + iree_string_builder_t* builder) { + iree_hal_null_driver_t* driver = iree_hal_null_driver_cast(base_driver); + + // TODO(null): add useful user-level information to the string builder for the + // given device_id. This is used by the tools in features like + // `--dump_devices` or may be used by hosting applications for diagnostics. + (void)driver; + + return iree_ok_status(); +} + +static iree_status_t iree_hal_null_driver_create_device_by_id( + iree_hal_driver_t* base_driver, iree_hal_device_id_t device_id, + iree_host_size_t param_count, const iree_string_pair_t* params, + iree_allocator_t host_allocator, iree_hal_device_t** out_device) { + iree_hal_null_driver_t* driver = iree_hal_null_driver_cast(base_driver); + + // TODO(null): use the provided params to overwrite the default options. The + // format of the params is implementation-defined. The params strings can be + // directly referenced if needed as the device creation is only allowed to + // access them during the create call below. + iree_hal_null_device_options_t options = + driver->options.default_device_options; + + // TODO(null): implement creation by device_id; this is mostly used as + // query_available_devices->create_device_by_id to avoid needing to expose + // device paths (which may not always be 1:1). This skeleton only has a single + // device so the ID is ignored. + (void)driver; + + return iree_hal_null_device_create(driver->identifier, &options, + host_allocator, out_device); +} + +static iree_status_t iree_hal_null_driver_create_device_by_path( + iree_hal_driver_t* base_driver, iree_string_view_t driver_name, + iree_string_view_t device_path, iree_host_size_t param_count, + const iree_string_pair_t* params, iree_allocator_t host_allocator, + iree_hal_device_t** out_device) { + iree_hal_null_driver_t* driver = iree_hal_null_driver_cast(base_driver); + + // TODO(null): use the provided params to overwrite the default options. The + // format of the params is implementation-defined. The params strings can be + // directly referenced if needed as the device creation is only allowed to + // access them during the create call below. + iree_hal_null_device_options_t options = + driver->options.default_device_options; + + // TODO(null): support parsing of the device_path. Note that a single driver + // may respond to multiple driver_name queries. Paths are + // implementation-specific and there may be multiple formats; for example, + // device UUID, PCI bus ID, ordinal as used by underlying APIs, etc. + (void)driver; + + return iree_hal_null_device_create(driver->identifier, &options, + host_allocator, out_device); +} + +static const iree_hal_driver_vtable_t iree_hal_null_driver_vtable = { + .destroy = iree_hal_null_driver_destroy, + .query_available_devices = iree_hal_null_driver_query_available_devices, + .dump_device_info = iree_hal_null_driver_dump_device_info, + .create_device_by_id = iree_hal_null_driver_create_device_by_id, + .create_device_by_path = iree_hal_null_driver_create_device_by_path, +}; diff --git a/runtime/src/iree/hal/drivers/null/driver.h b/runtime/src/iree/hal/drivers/null/driver.h new file mode 100644 index 000000000000..84b12c1beac8 --- /dev/null +++ b/runtime/src/iree/hal/drivers/null/driver.h @@ -0,0 +1,17 @@ +// Copyright 2024 The IREE Authors +// +// Licensed under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#ifndef IREE_HAL_DRIVERS_NULL_DRIVER_H_ +#define IREE_HAL_DRIVERS_NULL_DRIVER_H_ + +#include "iree/base/api.h" +#include "iree/hal/api.h" + +// NOTE: nothing in the skeleton implementation. Driver creation and adoption is +// part of the public API header. This header can contain internal types and +// functions. + +#endif // IREE_HAL_DRIVERS_NULL_DRIVER_H_ diff --git a/runtime/src/iree/hal/drivers/null/event.c b/runtime/src/iree/hal/drivers/null/event.c new file mode 100644 index 000000000000..5f1e413ca204 --- /dev/null +++ b/runtime/src/iree/hal/drivers/null/event.c @@ -0,0 +1,60 @@ +// Copyright 2024 The IREE Authors +// +// Licensed under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#include "iree/hal/drivers/null/event.h" + +typedef struct iree_hal_null_event_t { + iree_hal_resource_t resource; + iree_allocator_t host_allocator; +} iree_hal_null_event_t; + +static const iree_hal_event_vtable_t iree_hal_null_event_vtable; + +static iree_hal_null_event_t* iree_hal_null_event_cast( + iree_hal_event_t* base_value) { + IREE_HAL_ASSERT_TYPE(base_value, &iree_hal_null_event_vtable); + return (iree_hal_null_event_t*)base_value; +} + +iree_status_t iree_hal_null_event_create( + iree_hal_queue_affinity_t queue_affinity, iree_hal_event_flags_t flags, + iree_allocator_t host_allocator, iree_hal_event_t** out_event) { + IREE_ASSERT_ARGUMENT(out_event); + *out_event = NULL; + IREE_TRACE_ZONE_BEGIN(z0); + + iree_hal_null_event_t* event = NULL; + IREE_RETURN_AND_END_ZONE_IF_ERROR( + z0, + iree_allocator_malloc(host_allocator, sizeof(*event), (void**)&event)); + iree_hal_resource_initialize(&iree_hal_null_event_vtable, &event->resource); + event->host_allocator = host_allocator; + + // TODO(null): WIP API; this is a no-op today. + iree_status_t status = iree_ok_status(); + + if (iree_status_is_ok(status)) { + *out_event = (iree_hal_event_t*)event; + } else { + iree_hal_event_release((iree_hal_event_t*)event); + } + IREE_TRACE_ZONE_END(z0); + return status; +} + +static void iree_hal_null_event_destroy(iree_hal_event_t* base_event) { + iree_hal_null_event_t* event = iree_hal_null_event_cast(base_event); + iree_allocator_t host_allocator = event->host_allocator; + IREE_TRACE_ZONE_BEGIN(z0); + + iree_allocator_free(host_allocator, event); + + IREE_TRACE_ZONE_END(z0); +} + +static const iree_hal_event_vtable_t iree_hal_null_event_vtable = { + .destroy = iree_hal_null_event_destroy, +}; diff --git a/runtime/src/iree/hal/drivers/null/event.h b/runtime/src/iree/hal/drivers/null/event.h new file mode 100644 index 000000000000..68c11f44d1e6 --- /dev/null +++ b/runtime/src/iree/hal/drivers/null/event.h @@ -0,0 +1,18 @@ +// Copyright 2024 The IREE Authors +// +// Licensed under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#ifndef IREE_HAL_DRIVERS_NULL_EVENT_H_ +#define IREE_HAL_DRIVERS_NULL_EVENT_H_ + +#include "iree/base/api.h" +#include "iree/hal/api.h" + +// WIP API and may change. Mostly ignored for now. +iree_status_t iree_hal_null_event_create( + iree_hal_queue_affinity_t queue_affinity, iree_hal_event_flags_t flags, + iree_allocator_t host_allocator, iree_hal_event_t** out_event); + +#endif // IREE_HAL_DRIVERS_NULL_EVENT_H_ diff --git a/runtime/src/iree/hal/drivers/null/executable.c b/runtime/src/iree/hal/drivers/null/executable.c new file mode 100644 index 000000000000..a90d697d9d8d --- /dev/null +++ b/runtime/src/iree/hal/drivers/null/executable.c @@ -0,0 +1,86 @@ +// Copyright 2024 The IREE Authors +// +// Licensed under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#include "iree/hal/drivers/null/executable.h" + +typedef struct iree_hal_null_executable_t { + iree_hal_resource_t resource; + iree_allocator_t host_allocator; +} iree_hal_null_executable_t; + +static const iree_hal_executable_vtable_t iree_hal_null_executable_vtable; + +static iree_hal_null_executable_t* iree_hal_null_executable_cast( + iree_hal_executable_t* base_value) { + IREE_HAL_ASSERT_TYPE(base_value, &iree_hal_null_executable_vtable); + return (iree_hal_null_executable_t*)base_value; +} + +iree_status_t iree_hal_null_executable_create( + const iree_hal_executable_params_t* executable_params, + iree_allocator_t host_allocator, iree_hal_executable_t** out_executable) { + IREE_ASSERT_ARGUMENT(executable_params); + IREE_ASSERT_ARGUMENT(out_executable); + *out_executable = NULL; + IREE_TRACE_ZONE_BEGIN(z0); + + // Allocate storage for the executable and its associated data structures. + iree_hal_null_executable_t* executable = NULL; + IREE_RETURN_AND_END_ZONE_IF_ERROR( + z0, iree_allocator_malloc(host_allocator, sizeof(*executable), + (void**)&executable)); + iree_hal_resource_initialize(&iree_hal_null_executable_vtable, + &executable->resource); + executable->host_allocator = host_allocator; + + // TODO(null): load executable module(s). Note that the input data should be + // treated as untrusted and should be verified to the best ability the format + // provides. A target that cannot provide verification will be treated as + // unsafe. For JIT-style implementations as much work as possible should be + // done here so that errors can be propagated back to users - do not defer + // preparation. + // + // In general the executable should only retain information required to + // service the command buffer implementation that will be dispatching entry + // points within it. Optionally information can be retained for tracing and + // debugging. + // + // Implementations with flexible formats (ELF, etc) can directly use those for + // metadata as well with custom sections. If an implementation does not have a + // flexible format or support linking and requires several modules a wrapper + // can be used instead. In upstream IREE HALs Flatbuffers is used and is the + // preferred format (zero-copy, mmappable, verifiable, near header-only dep + // with no binary size or runtime overheads, etc) and is the easiest to use, + // but you do you. + iree_status_t status = + iree_make_status(IREE_STATUS_UNIMPLEMENTED, "executable not implemented"); + + if (iree_status_is_ok(status)) { + *out_executable = (iree_hal_executable_t*)executable; + } else { + iree_hal_executable_destroy((iree_hal_executable_t*)executable); + } + IREE_TRACE_ZONE_END(z0); + return status; +} + +static void iree_hal_null_executable_destroy( + iree_hal_executable_t* base_executable) { + iree_hal_null_executable_t* executable = + iree_hal_null_executable_cast(base_executable); + iree_allocator_t host_allocator = executable->host_allocator; + IREE_TRACE_ZONE_BEGIN(z0); + + // TODO(null): release any implementation resources. + + iree_allocator_free(host_allocator, executable); + + IREE_TRACE_ZONE_END(z0); +} + +static const iree_hal_executable_vtable_t iree_hal_null_executable_vtable = { + .destroy = iree_hal_null_executable_destroy, +}; diff --git a/runtime/src/iree/hal/drivers/null/executable.h b/runtime/src/iree/hal/drivers/null/executable.h new file mode 100644 index 000000000000..0107e1a14d4a --- /dev/null +++ b/runtime/src/iree/hal/drivers/null/executable.h @@ -0,0 +1,21 @@ +// Copyright 2024 The IREE Authors +// +// Licensed under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#ifndef IREE_HAL_DRIVERS_NULL_EXECUTABLE_H_ +#define IREE_HAL_DRIVERS_NULL_EXECUTABLE_H_ + +#include "iree/base/api.h" +#include "iree/hal/api.h" + +// Creates a {Null} executable from a binary in memory. Each executable may +// contain multiple entry points and be composed of several modules presented to +// the HAL as a single instance. See iree_hal_executable_params_t for more +// information about the lifetime of the resources referenced within. +iree_status_t iree_hal_null_executable_create( + const iree_hal_executable_params_t* executable_params, + iree_allocator_t host_allocator, iree_hal_executable_t** out_executable); + +#endif // IREE_HAL_DRIVERS_NULL_EXECUTABLE_H_ diff --git a/runtime/src/iree/hal/drivers/null/executable_cache.c b/runtime/src/iree/hal/drivers/null/executable_cache.c new file mode 100644 index 000000000000..d4f0ad6ad066 --- /dev/null +++ b/runtime/src/iree/hal/drivers/null/executable_cache.c @@ -0,0 +1,101 @@ +// Copyright 2024 The IREE Authors +// +// Licensed under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#include "iree/hal/drivers/null/executable_cache.h" + +#include "iree/hal/drivers/null/executable.h" + +typedef struct iree_hal_null_executable_cache_t { + iree_hal_resource_t resource; + iree_allocator_t host_allocator; +} iree_hal_null_executable_cache_t; + +static const iree_hal_executable_cache_vtable_t + iree_hal_null_executable_cache_vtable; + +static iree_hal_null_executable_cache_t* iree_hal_null_executable_cache_cast( + iree_hal_executable_cache_t* base_value) { + IREE_HAL_ASSERT_TYPE(base_value, &iree_hal_null_executable_cache_vtable); + return (iree_hal_null_executable_cache_t*)base_value; +} + +iree_status_t iree_hal_null_executable_cache_create( + iree_string_view_t identifier, iree_allocator_t host_allocator, + iree_hal_executable_cache_t** out_executable_cache) { + IREE_ASSERT_ARGUMENT(out_executable_cache); + *out_executable_cache = NULL; + IREE_TRACE_ZONE_BEGIN(z0); + + iree_hal_null_executable_cache_t* executable_cache = NULL; + IREE_RETURN_AND_END_ZONE_IF_ERROR( + z0, iree_allocator_malloc(host_allocator, sizeof(*executable_cache), + (void**)&executable_cache)); + iree_hal_resource_initialize(&iree_hal_null_executable_cache_vtable, + &executable_cache->resource); + executable_cache->host_allocator = host_allocator; + + // TODO(null): this default implementation is a no-op; if the implementation + // supports caching or has prohibitively expensive executable load times it is + // worth implementing this. Here any shared resources (compiler/JIT handles, + // device symbols, etc) can be retained and passed to all executables created + // from the cache. + iree_status_t status = iree_ok_status(); + + if (iree_status_is_ok(status)) { + *out_executable_cache = (iree_hal_executable_cache_t*)executable_cache; + } else { + iree_hal_executable_cache_release( + (iree_hal_executable_cache_t*)executable_cache); + } + IREE_TRACE_ZONE_END(z0); + return status; +} + +static void iree_hal_null_executable_cache_destroy( + iree_hal_executable_cache_t* base_executable_cache) { + iree_hal_null_executable_cache_t* executable_cache = + iree_hal_null_executable_cache_cast(base_executable_cache); + iree_allocator_t host_allocator = executable_cache->host_allocator; + IREE_TRACE_ZONE_BEGIN(z0); + + iree_allocator_free(host_allocator, executable_cache); + + IREE_TRACE_ZONE_END(z0); +} + +static bool iree_hal_null_executable_cache_can_prepare_format( + iree_hal_executable_cache_t* base_executable_cache, + iree_hal_executable_caching_mode_t caching_mode, + iree_string_view_t executable_format) { + // TODO(null): this query may be used with multiple executable_format args in + // cases where support is conditional (versions of a format) or multiple + // formats are handled. This matches the `IREE::HAL::ExecutableTargetAttr` + // format field as set in the compiler. + return iree_string_view_equal(executable_format, + iree_make_cstring_view("{null-executable}")); +} + +static iree_status_t iree_hal_null_executable_cache_prepare_executable( + iree_hal_executable_cache_t* base_executable_cache, + const iree_hal_executable_params_t* executable_params, + iree_hal_executable_t** out_executable) { + iree_hal_null_executable_cache_t* executable_cache = + iree_hal_null_executable_cache_cast(base_executable_cache); + + // TODO(null): add any extra args required. Usually this will be device + // symbols or handles that can be retained by the cache and passed to all + // executables created from it. + + return iree_hal_null_executable_create( + executable_params, executable_cache->host_allocator, out_executable); +} + +static const iree_hal_executable_cache_vtable_t + iree_hal_null_executable_cache_vtable = { + .destroy = iree_hal_null_executable_cache_destroy, + .can_prepare_format = iree_hal_null_executable_cache_can_prepare_format, + .prepare_executable = iree_hal_null_executable_cache_prepare_executable, +}; diff --git a/runtime/src/iree/hal/drivers/null/executable_cache.h b/runtime/src/iree/hal/drivers/null/executable_cache.h new file mode 100644 index 000000000000..519b8c05e18a --- /dev/null +++ b/runtime/src/iree/hal/drivers/null/executable_cache.h @@ -0,0 +1,23 @@ +// Copyright 2024 The IREE Authors +// +// Licensed under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#ifndef IREE_HAL_DRIVERS_NULL_EXECUTABLE_CACHE_H_ +#define IREE_HAL_DRIVERS_NULL_EXECUTABLE_CACHE_H_ + +#include "iree/base/api.h" +#include "iree/hal/api.h" + +// Creates a no-op executable cache that does not cache at all. +// This is useful to isolate pipeline caching behavior and verify compilation +// behavior. +// +// TODO(null): retain any shared resources (like device handles and symbols) +// that are needed to create executables. +iree_status_t iree_hal_null_executable_cache_create( + iree_string_view_t identifier, iree_allocator_t host_allocator, + iree_hal_executable_cache_t** out_executable_cache); + +#endif // IREE_HAL_DRIVERS_NULL_EXECUTABLE_CACHE_H_ diff --git a/runtime/src/iree/hal/drivers/null/registration/BUILD.bazel b/runtime/src/iree/hal/drivers/null/registration/BUILD.bazel new file mode 100644 index 000000000000..038eaf57e34b --- /dev/null +++ b/runtime/src/iree/hal/drivers/null/registration/BUILD.bazel @@ -0,0 +1,27 @@ +# Copyright 2024 The IREE Authors +# +# Licensed under the Apache License v2.0 with LLVM Exceptions. +# See https://llvm.org/LICENSE.txt for license information. +# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +load("//build_tools/bazel:build_defs.oss.bzl", "iree_runtime_cc_library") + +package( + default_visibility = ["//visibility:public"], + features = ["layering_check"], + licenses = ["notice"], # Apache 2.0 +) + +iree_runtime_cc_library( + name = "registration", + srcs = ["driver_module.c"], + hdrs = ["driver_module.h"], + defines = [ + "IREE_HAVE_HAL_NULL_DRIVER_MODULE=1", + ], + deps = [ + "//runtime/src/iree/base", + "//runtime/src/iree/hal", + "//runtime/src/iree/hal/drivers/null", + ], +) diff --git a/runtime/src/iree/hal/drivers/null/registration/CMakeLists.txt b/runtime/src/iree/hal/drivers/null/registration/CMakeLists.txt new file mode 100644 index 000000000000..4291d98018a0 --- /dev/null +++ b/runtime/src/iree/hal/drivers/null/registration/CMakeLists.txt @@ -0,0 +1,29 @@ +################################################################################ +# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from # +# runtime/src/iree/hal/drivers/null/registration/BUILD.bazel # +# # +# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary # +# CMake-only content. # +# # +# To disable autogeneration for this file entirely, delete this header. # +################################################################################ + +iree_add_all_subdirs() + +iree_cc_library( + NAME + registration + HDRS + "driver_module.h" + SRCS + "driver_module.c" + DEPS + iree::base + iree::hal + iree::hal::drivers::null + DEFINES + "IREE_HAVE_HAL_NULL_DRIVER_MODULE=1" + PUBLIC +) + +### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ### diff --git a/runtime/src/iree/hal/drivers/null/registration/driver_module.c b/runtime/src/iree/hal/drivers/null/registration/driver_module.c new file mode 100644 index 000000000000..7ffa9007f415 --- /dev/null +++ b/runtime/src/iree/hal/drivers/null/registration/driver_module.c @@ -0,0 +1,60 @@ +// Copyright 2024 The IREE Authors +// +// Licensed under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#include "iree/hal/drivers/null/registration/driver_module.h" + +#include "iree/base/api.h" +#include "iree/hal/drivers/null/api.h" + +static iree_status_t iree_hal_null_driver_factory_enumerate( + void* self, iree_host_size_t* out_driver_info_count, + const iree_hal_driver_info_t** out_driver_infos) { + // TODO(null): return multiple drivers if desired. This information must be + // static. The list here is just what is compiled into the binary and not + // expected to actually try to load or initialize drivers. + static const iree_hal_driver_info_t default_driver_info = { + .driver_name = IREE_SVL("null"), + .full_name = IREE_SVL("NULL Skeleton Driver"), + }; + *out_driver_info_count = 1; + *out_driver_infos = &default_driver_info; + return iree_ok_status(); +} + +static iree_status_t iree_hal_null_driver_factory_try_create( + void* self, iree_string_view_t driver_name, iree_allocator_t host_allocator, + iree_hal_driver_t** out_driver) { + // TODO(null): use your driver name - this will be the prefix when the user + // specifies the device (`--device=null://foo`). A single driver can support + // multiple prefixes if it wants. + if (!iree_string_view_equal(driver_name, IREE_SV("null"))) { + return iree_make_status(IREE_STATUS_UNAVAILABLE, + "no driver '%.*s' is provided by this factory", + (int)driver_name.size, driver_name.data); + } + + // TODO(null): populate options from flags. This driver module file is only + // used in native tools that have access to the flags library. Programmatic + // creation of the driver and devices will bypass this file and pass the + // options via this struct or key-value string parameters. + iree_hal_null_driver_options_t options; + iree_hal_null_driver_options_initialize(&options); + + iree_status_t status = iree_hal_null_driver_create( + driver_name, &options, host_allocator, out_driver); + + return status; +} + +IREE_API_EXPORT iree_status_t +iree_hal_null_driver_module_register(iree_hal_driver_registry_t* registry) { + static const iree_hal_driver_factory_t factory = { + .self = NULL, + .enumerate = iree_hal_null_driver_factory_enumerate, + .try_create = iree_hal_null_driver_factory_try_create, + }; + return iree_hal_driver_registry_register_factory(registry, &factory); +} diff --git a/runtime/src/iree/hal/drivers/null/registration/driver_module.h b/runtime/src/iree/hal/drivers/null/registration/driver_module.h new file mode 100644 index 000000000000..4ad0bd44f53a --- /dev/null +++ b/runtime/src/iree/hal/drivers/null/registration/driver_module.h @@ -0,0 +1,24 @@ +// Copyright 2024 The IREE Authors +// +// Licensed under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#ifndef IREE_HAL_DRIVER_NULL_REGISTRATION_DRIVER_MODULE_H_ +#define IREE_HAL_DRIVER_NULL_REGISTRATION_DRIVER_MODULE_H_ + +#include "iree/base/api.h" +#include "iree/hal/api.h" + +#ifdef __cplusplus +extern "C" { +#endif // __cplusplus + +IREE_API_EXPORT iree_status_t +iree_hal_null_driver_module_register(iree_hal_driver_registry_t* registry); + +#ifdef __cplusplus +} // extern "C" +#endif // __cplusplus + +#endif // IREE_HAL_DRIVER_NULL_REGISTRATION_DRIVER_MODULE_H_ diff --git a/runtime/src/iree/hal/drivers/null/semaphore.c b/runtime/src/iree/hal/drivers/null/semaphore.c new file mode 100644 index 000000000000..25ec7dc99fbb --- /dev/null +++ b/runtime/src/iree/hal/drivers/null/semaphore.c @@ -0,0 +1,172 @@ +// Copyright 2024 The IREE Authors +// +// Licensed under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#include "iree/hal/drivers/null/semaphore.h" + +#include "iree/hal/utils/semaphore_base.h" + +//===----------------------------------------------------------------------===// +// iree_hal_null_semaphore_t +//===----------------------------------------------------------------------===// + +typedef struct iree_hal_null_semaphore_t { + iree_hal_semaphore_t base; + iree_allocator_t host_allocator; +} iree_hal_null_semaphore_t; + +static const iree_hal_semaphore_vtable_t iree_hal_null_semaphore_vtable; + +static iree_hal_null_semaphore_t* iree_hal_null_semaphore_cast( + iree_hal_semaphore_t* base_value) { + IREE_HAL_ASSERT_TYPE(base_value, &iree_hal_null_semaphore_vtable); + return (iree_hal_null_semaphore_t*)base_value; +} + +iree_status_t iree_hal_null_semaphore_create( + uint64_t initial_value, iree_hal_semaphore_flags_t flags, + iree_allocator_t host_allocator, iree_hal_semaphore_t** out_semaphore) { + IREE_ASSERT_ARGUMENT(out_semaphore); + *out_semaphore = NULL; + IREE_TRACE_ZONE_BEGIN(z0); + + iree_hal_null_semaphore_t* semaphore = NULL; + IREE_RETURN_AND_END_ZONE_IF_ERROR( + z0, iree_allocator_malloc(host_allocator, sizeof(*semaphore), + (void**)&semaphore)); + iree_hal_semaphore_initialize(&iree_hal_null_semaphore_vtable, + &semaphore->base); + semaphore->host_allocator = host_allocator; + + // TODO(null): implement semaphores. Note that there is some basic support + // provided for timepoints as part of iree/hal/utils/semaphore_base.h but the + // actual synchronization aspects are handled by the implementation. + iree_status_t status = + iree_make_status(IREE_STATUS_UNIMPLEMENTED, "semaphore not implemented"); + + if (iree_status_is_ok(status)) { + *out_semaphore = &semaphore->base; + } else { + iree_hal_semaphore_release(&semaphore->base); + } + IREE_TRACE_ZONE_END(z0); + return status; +} + +static void iree_hal_null_semaphore_destroy( + iree_hal_semaphore_t* base_semaphore) { + iree_hal_null_semaphore_t* semaphore = + iree_hal_null_semaphore_cast(base_semaphore); + iree_allocator_t host_allocator = semaphore->host_allocator; + IREE_TRACE_ZONE_BEGIN(z0); + + iree_hal_semaphore_deinitialize(&semaphore->base); + iree_allocator_free(host_allocator, semaphore); + + IREE_TRACE_ZONE_END(z0); +} + +static iree_status_t iree_hal_null_semaphore_query( + iree_hal_semaphore_t* base_semaphore, uint64_t* out_value) { + *out_value = 0; + iree_hal_null_semaphore_t* semaphore = + iree_hal_null_semaphore_cast(base_semaphore); + + // TODO(null): return the current value of the semaphore by (depending on the + // implementation) making a syscall to get it. It's expected that the value + // may immediately change after being queried here. + + // TODO(null): if the value is IREE_HAL_SEMAPHORE_FAILURE_VALUE then return + // the failure status cached from the fail call by cloning it (like `return + // iree_status_clone(semaphore->failure_status)`). + + (void)semaphore; + iree_status_t status = iree_make_status(IREE_STATUS_UNIMPLEMENTED, + "semaphore query not implemented"); + + return status; +} + +static iree_status_t iree_hal_null_semaphore_signal( + iree_hal_semaphore_t* base_semaphore, uint64_t new_value) { + iree_hal_null_semaphore_t* semaphore = + iree_hal_null_semaphore_cast(base_semaphore); + + // TODO(null): validation is optional but encouraged if cheap: semaphores + // must always be signaled to a value that is greater than the previous value + // (not less-than-or-equal). + + // TODO(null): signals when the semaphore have failed should also fail and + // because failed semaphores have their value set to + // IREE_HAL_SEMAPHORE_FAILURE_VALUE that should happen naturally during + // validation. If not then an IREE_STATUS_DATA_LOSS or IREE_STATUS_ABORTED + // depending on how fatal such an occurrence is in the implementation. + // Data-loss usually indicates an abort()-worthy situation where graceful + // handling is not possible while Aborted indicates that an individual work + // stream may be invalid but unrelated work streams may still progress. + + (void)semaphore; + iree_status_t status = iree_make_status(IREE_STATUS_UNIMPLEMENTED, + "semaphore signal not implemented"); + + return status; +} + +static void iree_hal_null_semaphore_fail(iree_hal_semaphore_t* base_semaphore, + iree_status_t status) { + iree_hal_null_semaphore_t* semaphore = + iree_hal_null_semaphore_cast(base_semaphore); + const iree_status_code_t status_code = iree_status_code(status); + + // TODO(null): if the semaphore has already failed and has a status set then + // `IREE_IGNORE_ERROR(status)` and return without modifying anything. Note + // that it's possible for fail to be called concurrently from multiple + // threads. + + // TODO(null): set the value to `IREE_HAL_SEMAPHORE_FAILURE_VALUE` as expected + // by the API. + + // TODO(null): take ownership of the status (no need to clone, the caller is + // giving it to us) and keep it until the semaphore is destroyed. + + (void)semaphore; + (void)status_code; +} + +static iree_status_t iree_hal_null_semaphore_wait( + iree_hal_semaphore_t* base_semaphore, uint64_t value, + iree_timeout_t timeout) { + iree_hal_null_semaphore_t* semaphore = + iree_hal_null_semaphore_cast(base_semaphore); + + // TODO(null): if a failure status is set return + // `iree_status_from_code(IREE_STATUS_ABORTED)`. Avoid a full status as it may + // capture a backtrace and allocate and callers are expected to follow up a + // failed wait with a query to get the status. + + // TODO(null): prefer having a fast-path for if the semaphore is + // known-signaled in user-mode. This can usually avoid syscalls/ioctls and + // potential context switches in polling cases. + + // TODO(null): check for `iree_timeout_is_immediate(timeout)` and return + // immediately if the condition is not satisfied before waiting with + // `iree_status_from_code(IREE_STATUS_DEADLINE_EXCEEDED)`. Prefer the raw code + // status instead of a full status object as immediate timeouts are used when + // polling and a full status may capture a backtrace and allocate. + + (void)semaphore; + iree_status_t status = iree_make_status(IREE_STATUS_UNIMPLEMENTED, + "semaphore wait not implemented"); + + return status; +} + +static const iree_hal_semaphore_vtable_t iree_hal_null_semaphore_vtable = { + .destroy = iree_hal_null_semaphore_destroy, + .query = iree_hal_null_semaphore_query, + .signal = iree_hal_null_semaphore_signal, + .fail = iree_hal_null_semaphore_fail, + .wait = iree_hal_null_semaphore_wait, +}; diff --git a/runtime/src/iree/hal/drivers/null/semaphore.h b/runtime/src/iree/hal/drivers/null/semaphore.h new file mode 100644 index 000000000000..b27d361165a1 --- /dev/null +++ b/runtime/src/iree/hal/drivers/null/semaphore.h @@ -0,0 +1,23 @@ +// Copyright 2024 The IREE Authors +// +// Licensed under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#ifndef IREE_HAL_DRIVERS_NULL_SEMAPHORE_H_ +#define IREE_HAL_DRIVERS_NULL_SEMAPHORE_H_ + +#include "iree/base/api.h" +#include "iree/hal/api.h" + +//===----------------------------------------------------------------------===// +// iree_hal_null_semaphore_t +//===----------------------------------------------------------------------===// + +// Creates a {Null} semaphore used for ordering queue operations and +// synchronizing between host/device and device/device. +iree_status_t iree_hal_null_semaphore_create( + uint64_t initial_value, iree_hal_semaphore_flags_t flags, + iree_allocator_t host_allocator, iree_hal_semaphore_t** out_semaphore); + +#endif // IREE_HAL_DRIVERS_NULL_SEMAPHORE_H_