From d1d03cb7e620564fe81d4db70233f0c72f3de777 Mon Sep 17 00:00:00 2001 From: Lei Zhang Date: Mon, 7 Aug 2023 13:52:56 -0400 Subject: [PATCH] [metal] Move to hal/drivers and default build for Apple silicon (#14129) This commit moves the metal HAL driver from the `experimental/` directory to the `iree/hal/drivers/`, and turns on building by default for arm64 Apple devices. x86_64 Apple devices are not supported right now with missing features and such; that's to come later. These existing code is already extensively reviewed and passing various in-tree tests. Further improvements are for performance and tracked in https://github.com/openxla/iree/issues/14050. --- .github/CODEOWNERS | 1 + CMakeLists.txt | 16 ++-- docs/website/docs/index.md | 2 +- experimental/metal/tests/CMakeLists.txt | 7 -- .../metal/tests/stablehlo_ops/CMakeLists.txt | 73 ------------------- .../metal/tests/tosa_ops/CMakeLists.txt | 59 --------------- runtime/src/iree/hal/drivers/CMakeLists.txt | 4 + runtime/src/iree/hal/drivers/init.c | 9 +++ .../iree/hal/drivers}/metal/CMakeLists.txt | 6 +- .../src/iree/hal/drivers}/metal/README.md | 11 +-- .../src/iree/hal/drivers}/metal/api.h | 6 +- .../hal/drivers}/metal/builtin/CMakeLists.txt | 0 .../metal/builtin/copy_buffer_generic.metal | 0 .../metal/builtin/fill_buffer_generic.metal | 0 .../hal/drivers}/metal/builtin_executables.h | 8 +- .../hal/drivers}/metal/builtin_executables.m | 4 +- .../hal/drivers}/metal/cts/CMakeLists.txt | 4 +- .../hal/drivers}/metal/direct_allocator.h | 8 +- .../hal/drivers}/metal/direct_allocator.m | 4 +- .../drivers}/metal/direct_command_buffer.h | 12 +-- .../drivers}/metal/direct_command_buffer.m | 14 ++-- .../iree/hal/drivers}/metal/kernel_library.h | 6 +- .../iree/hal/drivers}/metal/kernel_library.m | 2 +- .../iree/hal/drivers}/metal/metal_buffer.h | 6 +- .../iree/hal/drivers}/metal/metal_buffer.m | 6 +- .../iree/hal/drivers}/metal/metal_device.h | 8 +- .../iree/hal/drivers}/metal/metal_device.m | 20 ++--- .../iree/hal/drivers}/metal/metal_driver.m | 4 +- .../hal/drivers}/metal/nop_executable_cache.h | 6 +- .../hal/drivers}/metal/nop_executable_cache.m | 5 +- .../iree/hal/drivers}/metal/pipeline_layout.h | 6 +- .../iree/hal/drivers}/metal/pipeline_layout.m | 2 +- .../metal/registration/CMakeLists.txt | 2 +- .../metal/registration/driver_module.c | 4 +- .../metal/registration/driver_module.h | 6 +- .../iree/hal/drivers}/metal/shared_event.h | 6 +- .../iree/hal/drivers}/metal/shared_event.m | 2 +- .../iree/hal/drivers}/metal/staging_buffer.h | 6 +- .../iree/hal/drivers}/metal/staging_buffer.m | 2 +- tests/e2e/stablehlo_ops/CMakeLists.txt | 71 ++++++++++++++++++ tests/e2e/tosa_ops/CMakeLists.txt | 53 ++++++++++++++ tools/build_config_template.txt.in | 1 + 42 files changed, 234 insertions(+), 238 deletions(-) delete mode 100644 experimental/metal/tests/CMakeLists.txt delete mode 100644 experimental/metal/tests/stablehlo_ops/CMakeLists.txt delete mode 100644 experimental/metal/tests/tosa_ops/CMakeLists.txt rename {experimental => runtime/src/iree/hal/drivers}/metal/CMakeLists.txt (84%) rename {experimental => runtime/src/iree/hal/drivers}/metal/README.md (97%) rename {experimental => runtime/src/iree/hal/drivers}/metal/api.h (97%) rename {experimental => runtime/src/iree/hal/drivers}/metal/builtin/CMakeLists.txt (100%) rename {experimental => runtime/src/iree/hal/drivers}/metal/builtin/copy_buffer_generic.metal (100%) rename {experimental => runtime/src/iree/hal/drivers}/metal/builtin/fill_buffer_generic.metal (100%) rename {experimental => runtime/src/iree/hal/drivers}/metal/builtin_executables.h (91%) rename {experimental => runtime/src/iree/hal/drivers}/metal/builtin_executables.m (98%) rename {experimental => runtime/src/iree/hal/drivers}/metal/cts/CMakeLists.txt (83%) rename {experimental => runtime/src/iree/hal/drivers}/metal/direct_allocator.h (87%) rename {experimental => runtime/src/iree/hal/drivers}/metal/direct_allocator.m (99%) rename {experimental => runtime/src/iree/hal/drivers}/metal/direct_command_buffer.h (86%) rename {experimental => runtime/src/iree/hal/drivers}/metal/direct_command_buffer.m (99%) rename {experimental => runtime/src/iree/hal/drivers}/metal/kernel_library.h (93%) rename {experimental => runtime/src/iree/hal/drivers}/metal/kernel_library.m (99%) rename {experimental => runtime/src/iree/hal/drivers}/metal/metal_buffer.h (89%) rename {experimental => runtime/src/iree/hal/drivers}/metal/metal_buffer.m (98%) rename {experimental => runtime/src/iree/hal/drivers}/metal/metal_device.h (84%) rename {experimental => runtime/src/iree/hal/drivers}/metal/metal_device.m (98%) rename {experimental => runtime/src/iree/hal/drivers}/metal/metal_driver.m (99%) rename {experimental => runtime/src/iree/hal/drivers}/metal/nop_executable_cache.h (82%) rename {experimental => runtime/src/iree/hal/drivers}/metal/nop_executable_cache.m (96%) rename {experimental => runtime/src/iree/hal/drivers}/metal/pipeline_layout.h (96%) rename {experimental => runtime/src/iree/hal/drivers}/metal/pipeline_layout.m (99%) rename {experimental => runtime/src/iree/hal/drivers}/metal/registration/CMakeLists.txt (93%) rename {experimental => runtime/src/iree/hal/drivers}/metal/registration/driver_module.c (96%) rename {experimental => runtime/src/iree/hal/drivers}/metal/registration/driver_module.h (72%) rename {experimental => runtime/src/iree/hal/drivers}/metal/shared_event.h (89%) rename {experimental => runtime/src/iree/hal/drivers}/metal/shared_event.m (99%) rename {experimental => runtime/src/iree/hal/drivers}/metal/staging_buffer.h (96%) rename {experimental => runtime/src/iree/hal/drivers}/metal/staging_buffer.m (98%) diff --git a/.github/CODEOWNERS b/.github/CODEOWNERS index d82e8e9de72b..96db1034d2fc 100644 --- a/.github/CODEOWNERS +++ b/.github/CODEOWNERS @@ -75,4 +75,5 @@ # Runtime /runtime/src/iree/ @benvanik /runtime/src/iree/hal/cts/ @ScottTodd +/runtime/src/iree/hal/drivers/metal/ @antiagainst /runtime/src/iree/hal/drivers/vulkan/ @antiagainst @ScottTodd diff --git a/CMakeLists.txt b/CMakeLists.txt index 9493b1b12079..441769753211 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -231,10 +231,18 @@ 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. +if(NOT APPLE OR NOT ${CMAKE_SYSTEM_PROCESSOR} MATCHES "arm64") + set(IREE_HAL_DRIVER_METAL_DEFAULT OFF) +endif() + option(IREE_HAL_DRIVER_CUDA "Enables the 'cuda' runtime HAL driver" ${IREE_HAL_DRIVER_CUDA_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_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}) @@ -295,6 +303,9 @@ endif() if(IREE_HAL_DRIVER_VULKAN) message(STATUS " - vulkan") endif() +if(IREE_HAL_DRIVER_METAL) + message(STATUS " - metal") +endif() if(IREE_EXTERNAL_HAL_DRIVERS) message(STATUS " + external: ${IREE_EXTERNAL_HAL_DRIVERS}") endif() @@ -336,11 +347,6 @@ set(IREE_EXTERNAL_ROCM_HAL_DRIVER_BINARY_DIR "${CMAKE_CURRENT_BINARY_DIR}/experi set(IREE_EXTERNAL_ROCM_HAL_DRIVER_TARGET "iree::experimental::rocm::registration") set(IREE_EXTERNAL_ROCM_HAL_DRIVER_REGISTER "iree_hal_rocm_driver_module_register") -set(IREE_EXTERNAL_METAL_HAL_DRIVER_SOURCE_DIR "${CMAKE_CURRENT_SOURCE_DIR}/experimental/metal") -set(IREE_EXTERNAL_METAL_HAL_DRIVER_BINARY_DIR "${CMAKE_CURRENT_BINARY_DIR}/experimental/metal") -set(IREE_EXTERNAL_METAL_HAL_DRIVER_TARGET "iree::experimental::metal::registration") -set(IREE_EXTERNAL_METAL_HAL_DRIVER_REGISTER "iree_hal_metal_driver_module_register") - #------------------------------------------------------------------------------- # Experimental WebGPU HAL driver #------------------------------------------------------------------------------- diff --git a/docs/website/docs/index.md b/docs/website/docs/index.md index a474b170bf19..f345f77202d9 100644 --- a/docs/website/docs/index.md +++ b/docs/website/docs/index.md @@ -53,7 +53,7 @@ Support for hardware accelerators and APIs is also included: - [x] Vulkan - [x] CUDA -- [ ] Metal (planned) +- [x] Metal (for Apple silicon devices) - [ ] WebGPU (planned) ## Project architecture diff --git a/experimental/metal/tests/CMakeLists.txt b/experimental/metal/tests/CMakeLists.txt deleted file mode 100644 index 33551b576974..000000000000 --- a/experimental/metal/tests/CMakeLists.txt +++ /dev/null @@ -1,7 +0,0 @@ -# Copyright 2023 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 - -iree_add_all_subdirs() diff --git a/experimental/metal/tests/stablehlo_ops/CMakeLists.txt b/experimental/metal/tests/stablehlo_ops/CMakeLists.txt deleted file mode 100644 index e4e05d081731..000000000000 --- a/experimental/metal/tests/stablehlo_ops/CMakeLists.txt +++ /dev/null @@ -1,73 +0,0 @@ -# Copyright 2023 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 - -iree_check_single_backend_test_suite( - NAME - check_metal-spirv_metal - SRCS - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/abs.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/add.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/batch_norm_inference.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/bitcast_convert.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/broadcast.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/broadcast_add.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/broadcast_in_dim.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/clamp.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/compare.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/concatenate.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/constant.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/convert.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/convolution.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/cosine.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/divide.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/dot.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/dot_general.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/dynamic_slice.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/dynamic_update_slice.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/exponential.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/exponential_fp16.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/exponential_minus_one.mlir" - # "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/fft.mlir" # TODO(#9583) - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/finite.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/floor.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/gather.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/iota.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/log.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/log_plus_one.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/maximum.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/minimum.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/multiply.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/negate.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/pad.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/pow.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/reduce.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/reduce_window.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/remainder.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/reshape.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/reverse.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/rng_normal.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/rng_uniform.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/round.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/rsqrt.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/scatter.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/scatter_dynamic.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/select.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/sine.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/slice.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/sort.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/sqrt.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/subtract.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/tanh.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/torch_index_select.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/transpose.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/stablehlo_ops/while.mlir" - TARGET_BACKEND - "metal-spirv" - DRIVER - "metal" - COMPILER_FLAGS - "--iree-input-type=stablehlo" -) diff --git a/experimental/metal/tests/tosa_ops/CMakeLists.txt b/experimental/metal/tests/tosa_ops/CMakeLists.txt deleted file mode 100644 index 161a9fc5ce5d..000000000000 --- a/experimental/metal/tests/tosa_ops/CMakeLists.txt +++ /dev/null @@ -1,59 +0,0 @@ -# Copyright 2023 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 - -iree_check_single_backend_test_suite( - NAME - check_metal-spirv_metal - SRCS - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/abs.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/add.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/arithmetic_right_shift.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/bitwise_and.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/bitwise_or.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/bitwise_xor.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/ceil.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/clamp.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/clz.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/const.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/equal.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/exp.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/floor.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/fully_connected.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/gather.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/greater.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/greater_equal.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/if.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/log.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/logical_left_shift.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/logical_right_shift.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/logical_right_shift_16.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/matmul.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/max_pool.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/maximum.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/minimum.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/mul.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/mul_shift.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/negate.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/pad.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/reciprocal.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/reduce.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/reshape.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/rsqrt.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/select.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/sigmoid.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/sub.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/table.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/tanh.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/transpose.mlir" - "${IREE_SOURCE_DIR}/tests/e2e/tosa_ops/while.mlir" - TARGET_BACKEND - "metal-spirv" - DRIVER - "metal" - COMPILER_FLAGS - "--iree-input-type=tosa" -) - diff --git a/runtime/src/iree/hal/drivers/CMakeLists.txt b/runtime/src/iree/hal/drivers/CMakeLists.txt index 711b493e8290..ac9906f8ea5d 100644 --- a/runtime/src/iree/hal/drivers/CMakeLists.txt +++ b/runtime/src/iree/hal/drivers/CMakeLists.txt @@ -126,6 +126,10 @@ 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() iree_cc_library( NAME diff --git a/runtime/src/iree/hal/drivers/init.c b/runtime/src/iree/hal/drivers/init.c index f39140d1213d..e2d1f659f4b6 100644 --- a/runtime/src/iree/hal/drivers/init.c +++ b/runtime/src/iree/hal/drivers/init.c @@ -22,6 +22,10 @@ #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( @@ -57,6 +61,11 @@ iree_hal_register_all_available_drivers(iree_hal_driver_registry_t* registry) { 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 + IREE_RETURN_AND_END_ZONE_IF_ERROR( z0, iree_hal_register_external_drivers(registry)); diff --git a/experimental/metal/CMakeLists.txt b/runtime/src/iree/hal/drivers/metal/CMakeLists.txt similarity index 84% rename from experimental/metal/CMakeLists.txt rename to runtime/src/iree/hal/drivers/metal/CMakeLists.txt index 88b75030fd0e..384005dff4bf 100644 --- a/experimental/metal/CMakeLists.txt +++ b/runtime/src/iree/hal/drivers/metal/CMakeLists.txt @@ -4,10 +4,6 @@ # See https://llvm.org/LICENSE.txt for license information. # SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# Set the root for package namespacing to the current directory. -set(IREE_PACKAGE_ROOT_DIR "${CMAKE_CURRENT_LIST_DIR}") -set(IREE_PACKAGE_ROOT_PREFIX "iree/experimental/metal") - iree_add_all_subdirs() iree_cc_library( @@ -43,8 +39,8 @@ iree_cc_library( iree::base::internal iree::base::internal::arena iree::base::internal::flatcc::parsing - iree::experimental::metal::builtin iree::hal + iree::hal::drivers::metal::builtin iree::hal::utils::buffer_transfer iree::hal::utils::resource_set iree::schemas::metal_executable_def_c_fbs diff --git a/experimental/metal/README.md b/runtime/src/iree/hal/drivers/metal/README.md similarity index 97% rename from experimental/metal/README.md rename to runtime/src/iree/hal/drivers/metal/README.md index 5eecd362d393..4ec5abe6fb4a 100644 --- a/experimental/metal/README.md +++ b/runtime/src/iree/hal/drivers/metal/README.md @@ -1,8 +1,8 @@ # IREE Metal HAL Driver This document lists technical details regarding the Metal HAL driver. Note that -the Metal HAL driver is a work in progress and experimental; this document is -expected to be updated along the way. +the Metal HAL driver is a work in progress; this document is expected to be +updated along the way. IREE provides a [Hardware Abstraction Layer (HAL)][iree-hal] as a common interface to different compute accelerators. IREE HAL's design draws inspiration @@ -26,10 +26,6 @@ already dominant ([macOS][macos-version-share], [iOS][ios-version-share]). ### Programming Languages and Libraries -The Metal HAL driver lives under the [`experimental/metal/`][iree-metal] -directory right now. Once more stable, it will be graduated into the -`runtime/src/iree/hal/drivers/` directory like other HAL drivers. - The Metal framework only exposes Objective-C or Swift programming language APIs. Metal HAL driver needs to inherit from common HAL abstraction definitions, which are in C. To minimize dependency and binary size and increase performance, we @@ -313,7 +309,6 @@ with the current active `MTLComputeCommandEncoder`: [macos-version-share]: https://gs.statcounter.com/macos-version-market-share/desktop/worldwide [ios-version-share]: https://developer.apple.com/support/app-store/ [iree-hal]: https://github.com/openxla/iree/tree/main/runtime/src/iree/hal -[iree-metal]: https://github.com/openxla/iree/tree/main/experimental/metal [hal-allocator]: https://github.com/openxla/iree/blob/main/runtime/src/iree/hal/allocator.h [hal-buffer]: https://github.com/openxla/iree/blob/main/runtime/src/iree/hal/buffer.h [hal-command-buffer]: https://github.com/openxla/iree/blob/main/runtime/src/iree/hal/command_buffer.h @@ -329,7 +324,7 @@ with the current active `MTLComputeCommandEncoder`: [metal-kernel-library]: https://github.com/openxla/iree/tree/main/experimental/metal/kernel_library.h [metal-shared-event]: https://github.com/openxla/iree/tree/main/experimental/metal/shared_event.h [metal-spirv-target]: https://github.com/openxla/iree/tree/main/compiler/src/iree/compiler/Dialect/HAL/Target/MetalSPIRV -[metal-builtin-kernels]: https://github.com/openxla/iree/tree/main/experimental/metal/builtin/ +[metal-builtin-kernels]: https://github.com/openxla/iree/tree/main/runtime/src/iree/hal/drivers/metal/builtin/ [mtl-argument-buffer]: https://developer.apple.com/documentation/metal/buffers/about_argument_buffers?language=objc [mtl-argument-encoder]: https://developer.apple.com/documentation/metal/mtlargumentencoder?language=objc [mtl-buffer]: https://developer.apple.com/documentation/metal/mtlbuffer?language=objc diff --git a/experimental/metal/api.h b/runtime/src/iree/hal/drivers/metal/api.h similarity index 97% rename from experimental/metal/api.h rename to runtime/src/iree/hal/drivers/metal/api.h index b2c9728f2621..f8056c49aa7c 100644 --- a/experimental/metal/api.h +++ b/runtime/src/iree/hal/drivers/metal/api.h @@ -6,8 +6,8 @@ // See iree/base/api.h for documentation on the API conventions used. -#ifndef IREE_EXPERIMENTAL_METAL_API_H_ -#define IREE_EXPERIMENTAL_METAL_API_H_ +#ifndef IREE_HAL_DRIVERS_METAL_API_H_ +#define IREE_HAL_DRIVERS_METAL_API_H_ #include "iree/base/api.h" #include "iree/hal/api.h" @@ -98,4 +98,4 @@ IREE_API_EXPORT iree_status_t iree_hal_metal_driver_create( } // extern "C" #endif // __cplusplus -#endif // IREE_EXPERIMENTAL_METAL_API_H_ +#endif // IREE_HAL_DRIVERS_METAL_API_H_ diff --git a/experimental/metal/builtin/CMakeLists.txt b/runtime/src/iree/hal/drivers/metal/builtin/CMakeLists.txt similarity index 100% rename from experimental/metal/builtin/CMakeLists.txt rename to runtime/src/iree/hal/drivers/metal/builtin/CMakeLists.txt diff --git a/experimental/metal/builtin/copy_buffer_generic.metal b/runtime/src/iree/hal/drivers/metal/builtin/copy_buffer_generic.metal similarity index 100% rename from experimental/metal/builtin/copy_buffer_generic.metal rename to runtime/src/iree/hal/drivers/metal/builtin/copy_buffer_generic.metal diff --git a/experimental/metal/builtin/fill_buffer_generic.metal b/runtime/src/iree/hal/drivers/metal/builtin/fill_buffer_generic.metal similarity index 100% rename from experimental/metal/builtin/fill_buffer_generic.metal rename to runtime/src/iree/hal/drivers/metal/builtin/fill_buffer_generic.metal diff --git a/experimental/metal/builtin_executables.h b/runtime/src/iree/hal/drivers/metal/builtin_executables.h similarity index 91% rename from experimental/metal/builtin_executables.h rename to runtime/src/iree/hal/drivers/metal/builtin_executables.h index a50ba6f5f32e..08fc065e3811 100644 --- a/experimental/metal/builtin_executables.h +++ b/runtime/src/iree/hal/drivers/metal/builtin_executables.h @@ -4,14 +4,14 @@ // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -#ifndef IREE_EXPERIMENTAL_METAL_BUILTIN_EXECUTABLES_H_ -#define IREE_EXPERIMENTAL_METAL_BUILTIN_EXECUTABLES_H_ +#ifndef IREE_HAL_DRIVERS_METAL_BUILTIN_EXECUTABLES_H_ +#define IREE_HAL_DRIVERS_METAL_BUILTIN_EXECUTABLES_H_ #import -#include "experimental/metal/kernel_library.h" #include "iree/base/api.h" #include "iree/hal/api.h" +#include "iree/hal/drivers/metal/kernel_library.h" #ifdef __cplusplus extern "C" { @@ -63,4 +63,4 @@ iree_status_t iree_hal_metal_builtin_executable_copy_buffer( } // extern "C" #endif // __cplusplus -#endif // IREE_EXPERIMENTAL_METAL_BUILTIN_EXECUTABLES_H_ +#endif // IREE_HAL_DRIVERS_METAL_BUILTIN_EXECUTABLES_H_ diff --git a/experimental/metal/builtin_executables.m b/runtime/src/iree/hal/drivers/metal/builtin_executables.m similarity index 98% rename from experimental/metal/builtin_executables.m rename to runtime/src/iree/hal/drivers/metal/builtin_executables.m index e20ff0269f52..77912ab9afee 100644 --- a/experimental/metal/builtin_executables.m +++ b/runtime/src/iree/hal/drivers/metal/builtin_executables.m @@ -4,14 +4,14 @@ // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -#include "experimental/metal/builtin_executables.h" +#include "iree/hal/drivers/metal/builtin_executables.h" #include -#include "experimental/metal/builtin/metal_buffer_kernels.h" #include "iree/base/api.h" #include "iree/base/tracing.h" #include "iree/hal/api.h" +#include "iree/hal/drivers/metal/builtin/metal_buffer_kernels.h" typedef struct iree_hal_metal_builtin_executable_data_t { const char* entry_point; diff --git a/experimental/metal/cts/CMakeLists.txt b/runtime/src/iree/hal/drivers/metal/cts/CMakeLists.txt similarity index 83% rename from experimental/metal/cts/CMakeLists.txt rename to runtime/src/iree/hal/drivers/metal/cts/CMakeLists.txt index 543d8c3c6ff5..4d99591808ce 100644 --- a/experimental/metal/cts/CMakeLists.txt +++ b/runtime/src/iree/hal/drivers/metal/cts/CMakeLists.txt @@ -8,7 +8,7 @@ iree_hal_cts_test_suite( DRIVER_NAME metal DRIVER_REGISTRATION_HDR - "experimental/metal/registration/driver_module.h" + "runtime/src/iree/hal/drivers/metal/registration/driver_module.h" DRIVER_REGISTRATION_FN "iree_hal_metal_driver_module_register" COMPILER_TARGET_BACKEND @@ -16,7 +16,7 @@ iree_hal_cts_test_suite( EXECUTABLE_FORMAT "\"MTLE\"" DEPS - iree::experimental::metal::registration + iree::hal::drivers::metal::registration EXCLUDED_TESTS # HAL event is unimplemented for Metal right now. "event" diff --git a/experimental/metal/direct_allocator.h b/runtime/src/iree/hal/drivers/metal/direct_allocator.h similarity index 87% rename from experimental/metal/direct_allocator.h rename to runtime/src/iree/hal/drivers/metal/direct_allocator.h index ae39a1817899..bbd460916ce8 100644 --- a/experimental/metal/direct_allocator.h +++ b/runtime/src/iree/hal/drivers/metal/direct_allocator.h @@ -4,14 +4,14 @@ // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -#ifndef IREE_EXPERIMENTAL_METAL_DIRECT_ALLOCATOR_H_ -#define IREE_EXPERIMENTAL_METAL_DIRECT_ALLOCATOR_H_ +#ifndef IREE_HAL_DRIVERS_METAL_DIRECT_ALLOCATOR_H_ +#define IREE_HAL_DRIVERS_METAL_DIRECT_ALLOCATOR_H_ #import -#include "experimental/metal/api.h" #include "iree/base/api.h" #include "iree/hal/api.h" +#include "iree/hal/drivers/metal/api.h" #ifdef __cplusplus extern "C" { @@ -44,4 +44,4 @@ id iree_hal_metal_allocator_command_queue( } // extern "C" #endif // __cplusplus -#endif // IREE_EXPERIMENTAL_METAL_DIRECT_ALLOCATOR_H_ +#endif // IREE_HAL_DRIVERS_METAL_DIRECT_ALLOCATOR_H_ diff --git a/experimental/metal/direct_allocator.m b/runtime/src/iree/hal/drivers/metal/direct_allocator.m similarity index 99% rename from experimental/metal/direct_allocator.m rename to runtime/src/iree/hal/drivers/metal/direct_allocator.m index 99968996b101..0b3a84122ae1 100644 --- a/experimental/metal/direct_allocator.m +++ b/runtime/src/iree/hal/drivers/metal/direct_allocator.m @@ -4,15 +4,15 @@ // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -#include "experimental/metal/direct_allocator.h" +#include "iree/hal/drivers/metal/direct_allocator.h" #import -#include "experimental/metal/metal_buffer.h" #include "iree/base/api.h" #include "iree/base/target_platform.h" #include "iree/base/tracing.h" #include "iree/hal/api.h" +#include "iree/hal/drivers/metal/metal_buffer.h" #if IREE_TRACING_FEATURES & IREE_TRACING_FEATURE_ALLOCATION_TRACKING static const char* IREE_HAL_METAL_ALLOCATOR_ID = "METAL"; diff --git a/experimental/metal/direct_command_buffer.h b/runtime/src/iree/hal/drivers/metal/direct_command_buffer.h similarity index 86% rename from experimental/metal/direct_command_buffer.h rename to runtime/src/iree/hal/drivers/metal/direct_command_buffer.h index 7ab8a6e073a8..49e82a514498 100644 --- a/experimental/metal/direct_command_buffer.h +++ b/runtime/src/iree/hal/drivers/metal/direct_command_buffer.h @@ -4,16 +4,16 @@ // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -#ifndef IREE_EXPERIMENTAL_METAL_METAL_COMMAND_BUFFER_H_ -#define IREE_EXPERIMENTAL_METAL_METAL_COMMAND_BUFFER_H_ +#ifndef IREE_HAL_DRIVERS_METAL_METAL_COMMAND_BUFFER_H_ +#define IREE_HAL_DRIVERS_METAL_METAL_COMMAND_BUFFER_H_ #import -#include "experimental/metal/api.h" -#include "experimental/metal/builtin_executables.h" -#include "experimental/metal/staging_buffer.h" #include "iree/base/internal/arena.h" #include "iree/hal/api.h" +#include "iree/hal/drivers/metal/api.h" +#include "iree/hal/drivers/metal/builtin_executables.h" +#include "iree/hal/drivers/metal/staging_buffer.h" #ifdef __cplusplus extern "C" { @@ -61,4 +61,4 @@ id iree_hal_metal_direct_command_buffer_handle( } // extern "C" #endif // __cplusplus -#endif // IREE_EXPERIMENTAL_METAL_METAL_COMMAND_BUFFER_H_ +#endif // IREE_HAL_DRIVERS_METAL_METAL_COMMAND_BUFFER_H_ diff --git a/experimental/metal/direct_command_buffer.m b/runtime/src/iree/hal/drivers/metal/direct_command_buffer.m similarity index 99% rename from experimental/metal/direct_command_buffer.m rename to runtime/src/iree/hal/drivers/metal/direct_command_buffer.m index c3dfdf029abb..a833ed9da547 100644 --- a/experimental/metal/direct_command_buffer.m +++ b/runtime/src/iree/hal/drivers/metal/direct_command_buffer.m @@ -4,20 +4,20 @@ // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -#include "experimental/metal/direct_command_buffer.h" +#include "iree/hal/drivers/metal/direct_command_buffer.h" #import -#include "experimental/metal/builtin_executables.h" -#include "experimental/metal/kernel_library.h" -#include "experimental/metal/metal_buffer.h" -#include "experimental/metal/metal_device.h" -#include "experimental/metal/pipeline_layout.h" -#include "experimental/metal/staging_buffer.h" #include "iree/base/api.h" #include "iree/base/target_platform.h" #include "iree/base/tracing.h" #include "iree/hal/api.h" +#include "iree/hal/drivers/metal/builtin_executables.h" +#include "iree/hal/drivers/metal/kernel_library.h" +#include "iree/hal/drivers/metal/metal_buffer.h" +#include "iree/hal/drivers/metal/metal_device.h" +#include "iree/hal/drivers/metal/pipeline_layout.h" +#include "iree/hal/drivers/metal/staging_buffer.h" #include "iree/hal/utils/resource_set.h" //===------------------------------------------------------------------------------------------===// diff --git a/experimental/metal/kernel_library.h b/runtime/src/iree/hal/drivers/metal/kernel_library.h similarity index 93% rename from experimental/metal/kernel_library.h rename to runtime/src/iree/hal/drivers/metal/kernel_library.h index 9cb11e712397..aa7c95742b77 100644 --- a/experimental/metal/kernel_library.h +++ b/runtime/src/iree/hal/drivers/metal/kernel_library.h @@ -4,8 +4,8 @@ // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -#ifndef IREE_EXPERIMENTAL_METAL_METAL_KERNEL_LIBRARY_H_ -#define IREE_EXPERIMENTAL_METAL_METAL_KERNEL_LIBRARY_H_ +#ifndef IREE_HAL_DRIVERS_METAL_KERNEL_LIBRARY_H_ +#define IREE_HAL_DRIVERS_METAL_KERNEL_LIBRARY_H_ #import #include @@ -61,4 +61,4 @@ iree_status_t iree_hal_metal_compile_msl_and_create_pipeline_object( } // extern "C" #endif // __cplusplus -#endif // IREE_EXPERIMENTAL_METAL_METAL_KERNEL_LIBRARY_H_ +#endif // IREE_HAL_DRIVERS_METAL_KERNEL_LIBRARY_H_ diff --git a/experimental/metal/kernel_library.m b/runtime/src/iree/hal/drivers/metal/kernel_library.m similarity index 99% rename from experimental/metal/kernel_library.m rename to runtime/src/iree/hal/drivers/metal/kernel_library.m index 156b7d3bb7be..f759f5e69c3b 100644 --- a/experimental/metal/kernel_library.m +++ b/runtime/src/iree/hal/drivers/metal/kernel_library.m @@ -4,7 +4,7 @@ // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -#include "experimental/metal/kernel_library.h" +#include "iree/hal/drivers/metal/kernel_library.h" #include diff --git a/experimental/metal/metal_buffer.h b/runtime/src/iree/hal/drivers/metal/metal_buffer.h similarity index 89% rename from experimental/metal/metal_buffer.h rename to runtime/src/iree/hal/drivers/metal/metal_buffer.h index b7f0b6549f57..2325b7436e87 100644 --- a/experimental/metal/metal_buffer.h +++ b/runtime/src/iree/hal/drivers/metal/metal_buffer.h @@ -4,8 +4,8 @@ // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -#ifndef IREE_EXPERIMENTAL_METAL_METAL_BUFFER_H_ -#define IREE_EXPERIMENTAL_METAL_METAL_BUFFER_H_ +#ifndef IREE_HAL_DRIVERS_METAL_METAL_BUFFER_H_ +#define IREE_HAL_DRIVERS_METAL_METAL_BUFFER_H_ #import @@ -37,4 +37,4 @@ id iree_hal_metal_buffer_handle(const iree_hal_buffer_t* buffer); } // extern "C" #endif // __cplusplus -#endif // IREE_EXPERIMENTAL_METAL_METAL_BUFFER_H_ +#endif // IREE_HAL_DRIVERS_METAL_METAL_BUFFER_H_ diff --git a/experimental/metal/metal_buffer.m b/runtime/src/iree/hal/drivers/metal/metal_buffer.m similarity index 98% rename from experimental/metal/metal_buffer.m rename to runtime/src/iree/hal/drivers/metal/metal_buffer.m index 5fa909edc01a..9ec1a85942aa 100644 --- a/experimental/metal/metal_buffer.m +++ b/runtime/src/iree/hal/drivers/metal/metal_buffer.m @@ -4,16 +4,16 @@ // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -#include "experimental/metal/metal_buffer.h" +#include "iree/hal/drivers/metal/metal_buffer.h" #import -#include "experimental/metal/direct_allocator.h" -#include "experimental/metal/metal_device.h" #include "iree/base/api.h" #include "iree/base/target_platform.h" #include "iree/base/tracing.h" #include "iree/hal/api.h" +#include "iree/hal/drivers/metal/direct_allocator.h" +#include "iree/hal/drivers/metal/metal_device.h" typedef struct iree_hal_metal_buffer_t { iree_hal_buffer_t base; diff --git a/experimental/metal/metal_device.h b/runtime/src/iree/hal/drivers/metal/metal_device.h similarity index 84% rename from experimental/metal/metal_device.h rename to runtime/src/iree/hal/drivers/metal/metal_device.h index 3348913fadf1..0013733fcfe7 100644 --- a/experimental/metal/metal_device.h +++ b/runtime/src/iree/hal/drivers/metal/metal_device.h @@ -4,14 +4,14 @@ // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -#ifndef IREE_EXPERIMENTAL_METAL_METAL_DEVICE_H_ -#define IREE_EXPERIMENTAL_METAL_METAL_DEVICE_H_ +#ifndef IREE_HAL_DRIVERS_METAL_METAL_DEVICE_H_ +#define IREE_HAL_DRIVERS_METAL_METAL_DEVICE_H_ #import -#include "experimental/metal/api.h" #include "iree/base/api.h" #include "iree/hal/api.h" +#include "iree/hal/drivers/metal/api.h" #ifdef __cplusplus extern "C" { @@ -34,4 +34,4 @@ const iree_hal_metal_device_params_t* iree_hal_metal_device_params( } // extern "C" #endif // __cplusplus -#endif // IREE_EXPERIMENTAL_METAL_METAL_DEVICE_H_ +#endif // IREE_HAL_DRIVERS_METAL_METAL_DEVICE_H_ diff --git a/experimental/metal/metal_device.m b/runtime/src/iree/hal/drivers/metal/metal_device.m similarity index 98% rename from experimental/metal/metal_device.m rename to runtime/src/iree/hal/drivers/metal/metal_device.m index e01442ce8aa4..4deb54df7fb0 100644 --- a/experimental/metal/metal_device.m +++ b/runtime/src/iree/hal/drivers/metal/metal_device.m @@ -4,19 +4,19 @@ // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -#include "experimental/metal/metal_device.h" - -#include "experimental/metal/api.h" -#include "experimental/metal/builtin_executables.h" -#include "experimental/metal/direct_allocator.h" -#include "experimental/metal/direct_command_buffer.h" -#include "experimental/metal/nop_executable_cache.h" -#include "experimental/metal/pipeline_layout.h" -#include "experimental/metal/shared_event.h" -#include "experimental/metal/staging_buffer.h" +#include "iree/hal/drivers/metal/metal_device.h" + #include "iree/base/api.h" #include "iree/base/tracing.h" #include "iree/hal/api.h" +#include "iree/hal/drivers/metal/api.h" +#include "iree/hal/drivers/metal/builtin_executables.h" +#include "iree/hal/drivers/metal/direct_allocator.h" +#include "iree/hal/drivers/metal/direct_command_buffer.h" +#include "iree/hal/drivers/metal/nop_executable_cache.h" +#include "iree/hal/drivers/metal/pipeline_layout.h" +#include "iree/hal/drivers/metal/shared_event.h" +#include "iree/hal/drivers/metal/staging_buffer.h" #include "iree/hal/utils/buffer_transfer.h" #include "iree/hal/utils/resource_set.h" diff --git a/experimental/metal/metal_driver.m b/runtime/src/iree/hal/drivers/metal/metal_driver.m similarity index 99% rename from experimental/metal/metal_driver.m rename to runtime/src/iree/hal/drivers/metal/metal_driver.m index 9f7be100cd01..993b5ac366c6 100644 --- a/experimental/metal/metal_driver.m +++ b/runtime/src/iree/hal/drivers/metal/metal_driver.m @@ -6,12 +6,12 @@ #import -#include "experimental/metal/api.h" -#include "experimental/metal/metal_device.h" #include "iree/base/api.h" #include "iree/base/target_platform.h" #include "iree/base/tracing.h" #include "iree/hal/api.h" +#include "iree/hal/drivers/metal/api.h" +#include "iree/hal/drivers/metal/metal_device.h" // Maximum device path length we support. The path is always a 16 character hex string. #define IREE_HAL_METAL_MAX_DEVICE_PATH_LENGTH 32 diff --git a/experimental/metal/nop_executable_cache.h b/runtime/src/iree/hal/drivers/metal/nop_executable_cache.h similarity index 82% rename from experimental/metal/nop_executable_cache.h rename to runtime/src/iree/hal/drivers/metal/nop_executable_cache.h index a75856561aed..507b4c348662 100644 --- a/experimental/metal/nop_executable_cache.h +++ b/runtime/src/iree/hal/drivers/metal/nop_executable_cache.h @@ -4,8 +4,8 @@ // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -#ifndef IREE_EXPERIMENTAL_METAL_METAL_NOP_EXECUTABLE_CACHE_H_ -#define IREE_EXPERIMENTAL_METAL_METAL_NOP_EXECUTABLE_CACHE_H_ +#ifndef IREE_HAL_DRIVERS_METAL_METAL_NOP_EXECUTABLE_CACHE_H_ +#define IREE_HAL_DRIVERS_METAL_METAL_NOP_EXECUTABLE_CACHE_H_ #import @@ -31,4 +31,4 @@ iree_status_t iree_hal_metal_nop_executable_cache_create( } // extern "C" #endif // __cplusplus -#endif // IREE_EXPERIMENTAL_METAL_METAL_NOP_EXECUTABLE_CACHE_H_ +#endif // IREE_HAL_DRIVERS_METAL_METAL_NOP_EXECUTABLE_CACHE_H_ diff --git a/experimental/metal/nop_executable_cache.m b/runtime/src/iree/hal/drivers/metal/nop_executable_cache.m similarity index 96% rename from experimental/metal/nop_executable_cache.m rename to runtime/src/iree/hal/drivers/metal/nop_executable_cache.m index ace2c82a4e52..347ce7d72a7a 100644 --- a/experimental/metal/nop_executable_cache.m +++ b/runtime/src/iree/hal/drivers/metal/nop_executable_cache.m @@ -4,15 +4,14 @@ // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -#include "experimental/metal/nop_executable_cache.h" +#include "iree/hal/drivers/metal/nop_executable_cache.h" -#include #include #include -#include "experimental/metal/kernel_library.h" #include "iree/base/api.h" #include "iree/base/tracing.h" +#include "iree/hal/drivers/metal/kernel_library.h" typedef struct iree_hal_metal_nop_executable_cache_t { // Abstract resource used for injecting reference counting and vtable; must be at offset 0. diff --git a/experimental/metal/pipeline_layout.h b/runtime/src/iree/hal/drivers/metal/pipeline_layout.h similarity index 96% rename from experimental/metal/pipeline_layout.h rename to runtime/src/iree/hal/drivers/metal/pipeline_layout.h index 63ed66b5a2b7..b97225b4b9d8 100644 --- a/experimental/metal/pipeline_layout.h +++ b/runtime/src/iree/hal/drivers/metal/pipeline_layout.h @@ -4,8 +4,8 @@ // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -#ifndef IREE_EXPERIMENTAL_METAL_PIPELINE_LAYOUT_H_ -#define IREE_EXPERIMENTAL_METAL_PIPELINE_LAYOUT_H_ +#ifndef IREE_HAL_DRIVERS_METAL_PIPELINE_LAYOUT_H_ +#define IREE_HAL_DRIVERS_METAL_PIPELINE_LAYOUT_H_ #include "iree/base/api.h" #include "iree/hal/api.h" @@ -103,4 +103,4 @@ iree_host_size_t iree_hal_metal_pipeline_layout_push_constant_count( } // extern "C" #endif // __cplusplus -#endif // IREE_EXPERIMENTAL_METAL_PIPELINE_LAYOUT_H_ +#endif // IREE_HAL_DRIVERS_METAL_PIPELINE_LAYOUT_H_ diff --git a/experimental/metal/pipeline_layout.m b/runtime/src/iree/hal/drivers/metal/pipeline_layout.m similarity index 99% rename from experimental/metal/pipeline_layout.m rename to runtime/src/iree/hal/drivers/metal/pipeline_layout.m index 4a687ed19365..b7899bb901b0 100644 --- a/experimental/metal/pipeline_layout.m +++ b/runtime/src/iree/hal/drivers/metal/pipeline_layout.m @@ -4,7 +4,7 @@ // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -#include "experimental/metal/pipeline_layout.h" +#include "iree/hal/drivers/metal/pipeline_layout.h" #include diff --git a/experimental/metal/registration/CMakeLists.txt b/runtime/src/iree/hal/drivers/metal/registration/CMakeLists.txt similarity index 93% rename from experimental/metal/registration/CMakeLists.txt rename to runtime/src/iree/hal/drivers/metal/registration/CMakeLists.txt index b4691a0c7ed1..0280231a4426 100644 --- a/experimental/metal/registration/CMakeLists.txt +++ b/runtime/src/iree/hal/drivers/metal/registration/CMakeLists.txt @@ -14,8 +14,8 @@ iree_cc_library( DEPS iree::base iree::base::core_headers - iree::experimental::metal iree::hal + iree::hal::drivers::metal DEFINES "IREE_HAVE_HAL_METAL_DRIVER_MODULE=1" PUBLIC diff --git a/experimental/metal/registration/driver_module.c b/runtime/src/iree/hal/drivers/metal/registration/driver_module.c similarity index 96% rename from experimental/metal/registration/driver_module.c rename to runtime/src/iree/hal/drivers/metal/registration/driver_module.c index da3569804afd..d2e20ebe340d 100644 --- a/experimental/metal/registration/driver_module.c +++ b/runtime/src/iree/hal/drivers/metal/registration/driver_module.c @@ -4,16 +4,16 @@ // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -#include "experimental/metal/registration/driver_module.h" +#include "iree/hal/drivers/metal/registration/driver_module.h" #include #include -#include "experimental/metal/api.h" #include "iree/base/api.h" #include "iree/base/internal/flags.h" #include "iree/base/status.h" #include "iree/base/tracing.h" +#include "iree/hal/drivers/metal/api.h" IREE_FLAG(bool, metal_serial_command_dispatch, false, "Serializes all commands within command buffers as if there were " diff --git a/experimental/metal/registration/driver_module.h b/runtime/src/iree/hal/drivers/metal/registration/driver_module.h similarity index 72% rename from experimental/metal/registration/driver_module.h rename to runtime/src/iree/hal/drivers/metal/registration/driver_module.h index abdef86556cb..72acfde1c27e 100644 --- a/experimental/metal/registration/driver_module.h +++ b/runtime/src/iree/hal/drivers/metal/registration/driver_module.h @@ -4,8 +4,8 @@ // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -#ifndef IREE_EXPERIMENTAL_METAL_REGISTRATION_DRIVER_MODULE_H_ -#define IREE_EXPERIMENTAL_METAL_REGISTRATION_DRIVER_MODULE_H_ +#ifndef IREE_HAL_DRIVERS_METAL_REGISTRATION_DRIVER_MODULE_H_ +#define IREE_HAL_DRIVERS_METAL_REGISTRATION_DRIVER_MODULE_H_ #include "iree/base/api.h" #include "iree/hal/api.h" @@ -21,4 +21,4 @@ iree_hal_metal_driver_module_register(iree_hal_driver_registry_t* registry); } // extern "C" #endif // __cplusplus -#endif // IREE_EXPERIMENTAL_METAL_REGISTRATION_DRIVER_MODULE_H_ +#endif // IREE_HAL_DRIVERS_METAL_REGISTRATION_DRIVER_MODULE_H_ diff --git a/experimental/metal/shared_event.h b/runtime/src/iree/hal/drivers/metal/shared_event.h similarity index 89% rename from experimental/metal/shared_event.h rename to runtime/src/iree/hal/drivers/metal/shared_event.h index 7961dcef39d5..4e3b8274e130 100644 --- a/experimental/metal/shared_event.h +++ b/runtime/src/iree/hal/drivers/metal/shared_event.h @@ -4,8 +4,8 @@ // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -#ifndef IREE_EXPERIMENTAL_METAL_METAL_SHARED_EVENT_H_ -#define IREE_EXPERIMENTAL_METAL_METAL_SHARED_EVENT_H_ +#ifndef IREE_HAL_DRIVERS_METAL_METAL_SHARED_EVENT_H_ +#define IREE_HAL_DRIVERS_METAL_METAL_SHARED_EVENT_H_ #import @@ -45,4 +45,4 @@ iree_status_t iree_hal_metal_shared_event_multi_wait( } // extern "C" #endif // __cplusplus -#endif // IREE_EXPERIMENTAL_METAL_METAL_SHARED_EVENT_H_ +#endif // IREE_HAL_DRIVERS_METAL_METAL_SHARED_EVENT_H_ diff --git a/experimental/metal/shared_event.m b/runtime/src/iree/hal/drivers/metal/shared_event.m similarity index 99% rename from experimental/metal/shared_event.m rename to runtime/src/iree/hal/drivers/metal/shared_event.m index 796707f226ef..fee92a13d3c4 100644 --- a/experimental/metal/shared_event.m +++ b/runtime/src/iree/hal/drivers/metal/shared_event.m @@ -4,7 +4,7 @@ // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -#include "experimental/metal/shared_event.h" +#include "iree/hal/drivers/metal/shared_event.h" #import diff --git a/experimental/metal/staging_buffer.h b/runtime/src/iree/hal/drivers/metal/staging_buffer.h similarity index 96% rename from experimental/metal/staging_buffer.h rename to runtime/src/iree/hal/drivers/metal/staging_buffer.h index c46df94e2471..63c6ea1c3bd3 100644 --- a/experimental/metal/staging_buffer.h +++ b/runtime/src/iree/hal/drivers/metal/staging_buffer.h @@ -4,8 +4,8 @@ // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -#ifndef IREE_EXPERIMENTAL_METAL_STAGING_BUFFER_H_ -#define IREE_EXPERIMENTAL_METAL_STAGING_BUFFER_H_ +#ifndef IREE_HAL_DRIVERS_METAL_STAGING_BUFFER_H_ +#define IREE_HAL_DRIVERS_METAL_STAGING_BUFFER_H_ #import @@ -99,4 +99,4 @@ void iree_hal_metal_staging_buffer_decrease_refcount( } // extern "C" #endif // __cplusplus -#endif // IREE_EXPERIMENTAL_METAL_STAGING_BUFFER_H_ +#endif // IREE_HAL_DRIVERS_METAL_STAGING_BUFFER_H_ diff --git a/experimental/metal/staging_buffer.m b/runtime/src/iree/hal/drivers/metal/staging_buffer.m similarity index 98% rename from experimental/metal/staging_buffer.m rename to runtime/src/iree/hal/drivers/metal/staging_buffer.m index 08791b4678e2..a7eb80603236 100644 --- a/experimental/metal/staging_buffer.m +++ b/runtime/src/iree/hal/drivers/metal/staging_buffer.m @@ -4,7 +4,7 @@ // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -#include "experimental/metal/staging_buffer.h" +#include "iree/hal/drivers/metal/staging_buffer.h" #include diff --git a/tests/e2e/stablehlo_ops/CMakeLists.txt b/tests/e2e/stablehlo_ops/CMakeLists.txt index 451a00373dd5..ccd338895788 100644 --- a/tests/e2e/stablehlo_ops/CMakeLists.txt +++ b/tests/e2e/stablehlo_ops/CMakeLists.txt @@ -533,3 +533,74 @@ iree_check_single_backend_test_suite( "--iree-codegen-gpu-native-math-precision=true" # TODO(#11321): Infer/flip default ) +iree_check_single_backend_test_suite( + NAME + check_metal-spirv_metal + SRCS + "abs.mlir" + "add.mlir" + "batch_norm_inference.mlir" + "bitcast_convert.mlir" + "broadcast.mlir" + "broadcast_add.mlir" + "broadcast_in_dim.mlir" + "clamp.mlir" + "compare.mlir" + "complex.mlir" + "concatenate.mlir" + "constant.mlir" + "convert.mlir" + "convolution.mlir" + "cosine.mlir" + "divide.mlir" + "dot.mlir" + "dot_bf16.mlir" + "dot_general.mlir" + "dynamic_slice.mlir" + "dynamic_update_slice.mlir" + "exponential.mlir" + "exponential_fp16.mlir" + "exponential_minus_one.mlir" + "fft.mlir" + "finite.mlir" + "floor.mlir" + "gather.mlir" + "iota.mlir" + "log.mlir" + "log_plus_one.mlir" + "maximum.mlir" + "minimum.mlir" + "multiply.mlir" + "negate.mlir" + "pad.mlir" + "philox.mlir" + "pow.mlir" + "reduce.mlir" + "reduce_window.mlir" + "remainder.mlir" + "reshape.mlir" + "reverse.mlir" + "rng_normal.mlir" + "rng_uniform.mlir" + "round.mlir" + "rsqrt.mlir" + "scatter.mlir" + "scatter_dynamic.mlir" + "select.mlir" + "sine.mlir" + "slice.mlir" + "sort.mlir" + "sqrt.mlir" + "subtract.mlir" + "tanh.mlir" + "three_fry.mlir" + "torch_index_select.mlir" + "transpose.mlir" + "while.mlir" + TARGET_BACKEND + "metal-spirv" + DRIVER + "metal" + COMPILER_FLAGS + "--iree-input-type=stablehlo" +) diff --git a/tests/e2e/tosa_ops/CMakeLists.txt b/tests/e2e/tosa_ops/CMakeLists.txt index 296c223abe37..a4a0d7ea99ab 100644 --- a/tests/e2e/tosa_ops/CMakeLists.txt +++ b/tests/e2e/tosa_ops/CMakeLists.txt @@ -276,3 +276,56 @@ iree_check_single_backend_test_suite( "--iree-input-type=tosa" "--iree-codegen-gpu-native-math-precision=true" # TODO(#11321): Infer/flip default ) + +iree_check_single_backend_test_suite( + NAME + check_metal-spirv_metal + SRCS + "abs.mlir" + "add.mlir" + "arithmetic_right_shift.mlir" + "bitwise_and.mlir" + "bitwise_or.mlir" + "bitwise_xor.mlir" + "ceil.mlir" + "clamp.mlir" + "clz.mlir" + "const.mlir" + "equal.mlir" + "exp.mlir" + "floor.mlir" + "fully_connected.mlir" + "gather.mlir" + "greater.mlir" + "greater_equal.mlir" + "if.mlir" + "log.mlir" + "logical_left_shift.mlir" + "logical_right_shift.mlir" + "logical_right_shift_16.mlir" + "matmul.mlir" + "max_pool.mlir" + "maximum.mlir" + "minimum.mlir" + "mul.mlir" + "mul_shift.mlir" + "negate.mlir" + "pad.mlir" + "reciprocal.mlir" + "reduce.mlir" + "reshape.mlir" + "rsqrt.mlir" + "select.mlir" + "sigmoid.mlir" + "sub.mlir" + "table.mlir" + "tanh.mlir" + "transpose.mlir" + "while.mlir" + TARGET_BACKEND + "metal-spirv" + DRIVER + "metal" + COMPILER_FLAGS + "--iree-input-type=tosa" +) diff --git a/tools/build_config_template.txt.in b/tools/build_config_template.txt.in index a0fb76675285..fa4263e125cb 100644 --- a/tools/build_config_template.txt.in +++ b/tools/build_config_template.txt.in @@ -2,6 +2,7 @@ IREE_HAL_DRIVER_CUDA=${IREE_HAL_DRIVER_CUDA} IREE_HAL_DRIVER_LOCAL_SYNC=${IREE_HAL_DRIVER_LOCAL_SYNC} IREE_HAL_DRIVER_LOCAL_TASK=${IREE_HAL_DRIVER_LOCAL_TASK} IREE_HAL_DRIVER_VULKAN=${IREE_HAL_DRIVER_VULKAN} +IREE_HAL_DRIVER_METAL=${IREE_HAL_DRIVER_METAL} IREE_HAL_EXECUTABLE_LOADER_SYSTEM_LIBRARY=${IREE_HAL_EXECUTABLE_LOADER_SYSTEM_LIBRARY} IREE_HAL_EXECUTABLE_LOADER_EMBEDDED_ELF=${IREE_HAL_EXECUTABLE_LOADER_EMBEDDED_ELF} IREE_HAL_EXECUTABLE_LOADER_VMVX_MODULE=${IREE_HAL_EXECUTABLE_LOADER_VMVX_MODULE}