Removing Metal backend until it can be replaced.
This commit can be used to revive the configuration for the driver
and view a partially-ported version of it. The actual implementation
will need to be rebased onto the C type system and written in
Objective C. Note that the compiler backend has not been removed - we
can use that for verifying metal shader output.
Issue #4370 tracks reviving this.
diff --git a/CMakeLists.txt b/CMakeLists.txt
index 03a213b..287ccf5 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -115,7 +115,6 @@
# List of all HAL drivers to be built by default:
set(IREE_ALL_HAL_DRIVERS
DyLib
- Metal
VMLA
Vulkan
)
@@ -126,9 +125,6 @@
# For Apple platforms we need to use Metal instead of Vulkan.
if(APPLE)
list(REMOVE_ITEM IREE_HAL_DRIVERS_TO_BUILD Vulkan)
- else()
- # And Metal isn't available on non-Apple platforms for sure.
- list(REMOVE_ITEM IREE_HAL_DRIVERS_TO_BUILD Metal)
endif()
endif()
message(STATUS "Building HAL drivers: ${IREE_HAL_DRIVERS_TO_BUILD}")
diff --git a/docs/design_docs/metal_hal_driver.md b/docs/design_docs/metal_hal_driver.md
deleted file mode 100644
index 064e2bb..0000000
--- a/docs/design_docs/metal_hal_driver.md
+++ /dev/null
@@ -1,371 +0,0 @@
-# Metal HAL Driver
-
-This document lists technical details regarding the Metal HAL driver. Note that
-the Metal HAL driver is working 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
-from modern GPU architecture and APIs; so implementing a HAL driver using modern
-GPU APIs is generally straightforward. This applies to the Metal HAL driver.
-
-## Overall Design Choices
-
-### Metal Versions
-
-The Metal HAL driver expects Metal 2+. Metal 2 introduces useful features like
-argument buffer, performance shaders, and others, that can improve performance
-and make IREE HAL implementation simpler. Metal 2 was released late 2017 and are
-supported since macOS High Sierra and iOS 11. It is already dominant
-([macOS][macos-version-share], [iOS][ios-version-share]) right now.
-
-### Programming Languages and Libraries
-
-The Metal HAL driver lives under the [`iree/hal/metal/`][iree-metal] directory.
-Header (`.h`) and implementation (`.mm`) files are put adjacent to each other.
-
-The Metal framework only exposes Objective-C or Swift programming language APIs.
-Metal HAL driver needs to inherit from common HAL abstraction classes, which are
-C++. So we use [Objective-C++][objcxx] for implementing the Metal HAL driver.
-The headers try to stay with pure C/C++ syntax when possible, except for
-`#import <Metal/Metal.h>` and using Metal `id` types.
-
-### Object Lifetime Management
-
-Objective-C uses refcount for tracking object lifetime and managing memory. This
-is traditionally done manually by sending `retain` and `release` messages to
-Objective-C objects. Modern Objective-C allows developers to opt in to use
-[Automatic Reference Counting][objc-arc] to let the compiler to automatically
-deduce and insert `retain`/`release` where possible to simplify the burdern of
-manual management.
-
-We don't use ARC in the Metal HAL driver given that IREE has its own object
-[refcount][iree-refptr] and lifetime management mechanism. Metal HAL GPU objects
-are tracked with that to be consistent with others. Each Metal HAL GPU object
-`retain`s the underlying Metal `id<MTL*>` object on construction and `release`s
-on destruction.
-
-## GPU Objects
-
-Metal is one of the main modern GPU APIs that provide more explicit control over
-the hardware. The mapping between IREE HAL classes and Metal protocols are
-relatively straightforward:
-
-IREE HAL Class | Metal Protocol
-:----------------------------------------: | :------------:
-[`hal::Driver`][hal-driver] | N/A
-[`hal::Device`][hal-device] | [`MTLDevice`][mtl-device]
-[`hal::CommandQueue`][hal-command-queue] | [`MTLCommandQueue`][mtl-command-queue]
-[`hal::CommandBuffer`][hal-command-buffer] | [`MTLCommandBuffer`][mtl-command-buffer]
-[`hal::Semaphore`][hal-semaphore] | [`MTLSharedEvent`][mtl-shared-event]
-[`hal::Allocator`][hal-allocator] | N/A
-[`hal::Buffer`][hal-buffer] | [`MTLBuffer`][mtl-buffer]
-[`hal::Executable`][hal-executable] | [`MTLLibrary`][mtl-library]
-[`hal::ExecutableCache`][hal-executable-cache] | N/A
-[`hal::DescriptorSetLayout`][hal-descriptor-set-layout] | N/A
-[`hal::DescriptorSet`][hal-descriptor-set] | N/A
-[`hal::ExecutableLayout`][hal-executable-layout] | N/A
-
-In the following subsections, we go over each pair to provide more details.
-
-### Driver
-
-There is no native driver abstraction in Metal. IREE's Metal HAL driver still
-provides a [`hal::metal::MetalDriver`][metal-driver] subclass inheriting from
-common [`hal::Driver`][hal-driver] class. `hal::metal::MetalDriver` just
-`retain`s all available Metal devices in the system during its lifetime to
-provide similar interface as other HAL drivers.
-
-### Device
-
-[`hal::metal::MetalDevice`][metal-device] inherits [`hal::Device`][hal-device]
-to provide the interface to Metal GPU device by wrapping a `id<MTLDevice>`. Upon
-construction, `hal::metal::MetalDevice` creates and retains one queue for both
-dispatch and transfer during its lifetime.
-
-Metal requres command buffers to be created from a `MTLCommandQueue`. In IREE
-HAL, command buffers are directly created from the `hal::Device`.
-`hal::metal::MetalDevice` chooses the proper queue to create the command buffer
-under the hood.
-
-### Command queue
-
-IREE HAL command queue follows Vulkan for modelling submission. Specifically,
-`hal::CommandQueue::Submit()` takes a `SubmissionBatch`, which contains a list
-of waiting `hal::Semaphore`s, a list of command buffers, and a list signaling
-`hal::Semaphore`s. There is no direct mapping in Metal; so
-[`hal::metal::MetalCommandQueue`][metal-command-queue] performs the submission
-in three steps:
-
-1. Create a new `MTLCommandBuffer` to `encodeWaitForEvent:value` for all
- waiting `hal::Semaphore`s and commit this command buffer.
-1. Commit all command buffers in the `SubmissionBatch`.
-1. Create a new `MTLCommandBuffer` to `encodeSignalEvent:value` for all
- signaling `hal::Semaphore`s and commit this command buffer.
-
-There is also no direct `WaitIdle()` for
-[`MTLCommandQueue`][mtl-command-queue]s. `hal::metal::MetalCommandQueue`
-implements `WaitIdle()` by committing an empty `MTLCommandBuffer` and
-registering a complete handler for it to signal a semaphore to wake the current
-thread, which is put into sleep by waiting on the semaphore.
-
-### Command buffer
-
-In Metal, commands are recorded into a command buffer with three different kinds
-of [command encoders][mtl-command-encoder]: `MTLRenderCommandEncoder`,
-`MTLComputeCommandEncoder`, `MTLBlitCommandEncoder`, and
-`MTLParallelRenderCommandEncoder`. Each encoder has its own create/end call.
-There is no overall begin/end call for the whold command buffer. So even
-[`hal::metal::MetalCommandBuffer`][metal-command-buffer] implements an overall
-`Begin()`/`End()` call, under the hood it may create a new command encoder for a
-specific API call.
-
-### Timeline semaphore
-
-[`hal::Semaphore`][hal-semaphore] allows host->device, device->host, host->host,
-and device->device synchronization. It maps to Vulkan timeline semaphore. In
-Metal world, the counterpart would be [`MTLSharedEvent`][mtl-shared-event]. Most
-of the `hal::Semaphore` APIs are simple to implement in
-[`MetalSharedEvent`][metal-shared-event], with `Wait()` as an exception. A
-listener is registered on the `MTLSharedEvent` with
-`notifyListener:atValue:block:` to singal a semaphore to wake the current
-thread, which is put into sleep by waiting on the semaphore.
-
-### Allocator
-
-At the moment the Metal HAL driver just has a very simple
-[`hal::Allocator`][hal-allocator] implementation. It just wraps a `MTLDevice`
-and redirects all allocation requests to the `MTLDevice`. No page/pool/slab or
-whatever. This is only meant to get started. In the future we should have a
-better memory allocation library, probably by layering the
-[Vulkan Memory Allocator][vma] on top of [`MTLHeap`][mtl-heap].
-
-### Buffer
-
-IREE [`hal::Buffer`][hal-buffer] maps Metal `MTLBuffer`. See
-[Memory Management](#memory-management) for more details.
-
-### Executable
-
-IREE [`hal::Executable`][hal-executable] represents a GPU program archive with
-a driver-defined format. It maps naturally to Metal [`MTLLibrary`][mtl-library].
-An entry point in a `MTLLibrary` is a [`MTLFunction`][mtl-function]. We define
-[`hal::metal::MetalKernelLibrary`][metal-kernel-library] to wrap around a
-`MTLLibrary`, its `MTLFunction`s, and also `MTLComputePipelineState` objects
-constructed from `MTLFunction`s.
-
-### Executable cache
-
-IREE [`hal::ExecutableCache`][hal-executable-cache] is modelling a cache of
-preprared GPU executables for a particular device. At the moment the Metal
-HAL driver does not peforming any cache on GPU programs; it simply reads the
-program from the FlatBuffer and hands it over to Metal driver.
-
-### DescriptorSetLayout, DescriptorSet, ExecutableLayout
-
-See [Resource descriptors](#resource-descriptors) for more details.
-
-## Compute Pipeline
-
-### Shader/kernel compilation
-
-Metal has [Metal Shading Language (MSL)][msl-spec] for authoring graphics
-shaders and compute kernels. MSL source code can be directly consumed by the
-Metal framework at run-time; it can also be compiled first into an opaque
-library using [command-line tools][msl-cl-library] at build-time.
-
-IREE uses compilers to compile ML models expressed with high-level op semantics
-down to GPU native source format. This is also the case for the Metal HAL
-driver. Metal does not provide an open intermediate language; we reuse the
-[SPIR-V code generation pipeline][spirv-codegen] and then cross compile the
-generated SPIR-V into MSL source with [SPIRV-Cross][spirv-cross]. This is
-actually a fair common practice for targeting multiple GPU APIs in graphics
-programming world. For example, the Vulkan implmenation in macOS/iOs,
-[MoltenVK][moltenvk], is also doing the same for shaders/kernels. The path
-is actually quite robust, as demonstrated by various games on top of MoltenVK.
-
-Therefore, in IREE, we have a [`MetalSPIRVTargetBackend`][metal-spirv-target],
-which pulls in the normal MHLO to Linalg and Linalg to SPIR-V passes to form
-the compilation pipeline. The difference would be to provide a suitable
-SPIR-V target environment to drive the compilation, which one can derive from
-the Metal GPU families to target. (Not implemented yet; TODO for the future.)
-The serialization step differs from
-[`VulkanSPIRVTargetBackend`][vulkan-spirv-target] too: following the normal
-SPIR-V serialization step, we additionally need to invoke SPRIV-Cross to
-cross compile the generated SPIR-V into MSL, and then compile and/or serialize
-the MSL source/library.
-
-IREE uses [FlatBuffer][flatbuffer] to encode the whole workload module,
-including both GPU shader/kernel (called executable in IREE terminology) and
-CPU scheduling logic. The GPU executables are embedded as part of the module's
-FlatBuffer, which are [`mmap`][mmap]ped when IREE runs.
-
-For the Metal HAL driver, it means we need to embed the MSL kernels inside the
-module FlatBuffer. Right now we just encode the MSL source strings and compile
-them at Metal run-time. In the future this should be changed to allow encoding
-the library instead.
-
-### Workgroup/threadgroup size
-
-When dispatching a compute kernel in Metal, we need to specify the number of
-thread groups in grid and the number of threads in thread group. Both are
-3-D vectors. IREE HAL, which follows Vulkan, calls them workgroup count and
-workgroup size, respectively.
-
-In Vulkan programming model, workgroup count and workgroup size are specified at
-different places: the former is given when invoking
-[`vkCmdDispatch()`][vulkan-cmd-dispatch], while the later is encoded in the
-dispatched SPIR-V code. This split does not match the Metal model, where we
-specify both in the API with `dispatchThreads:threadsPerThreadgroup:`.
-
-As said in [shader/kernel compilation](#shader-kernel-compilation), MSL kernels
-are cross compiled from SPIR-V code and then embeded in the module FlatBuffer.
-The module FlatBuffer provides us a way to convey the threadgroup/workgroup size
-information extracted from the SPIR-V code. We encode an additional 3-D vector
-for each entry point and use it as the threadgroup size when later dispatching
-the `MTLFunction` corresponding to the entry point.
-
-### Resource descriptors
-
-A descriptor is an opaque handle pointing to a resource that is accessed in
-the compute kernel. IREE's HAL is inspired by the Vulkan API; it models several
-concepts related to GPU resource management explicitly:
-
-* [`hal::DescriptorSetLayout`][hal-descriptor-set-layout]: a schema for
- describing an array of descriptor bindings. Each descriptor binding specifies
- the resource type, access mode and other information.
-* [`hal::DescriptorSet`][hal-descriptor-set]: a concrete set of resources that
- gets bound to a compute pipeline in a batch. It must match the
- `DescriptorSetLayout` describing its layout. `DescriptorSet` can be thought as
- the "object" from the `DescriptorSetLayout` "class".
-* [`hal::ExecutableLayout`][hal-executable-layout]: a schema for describing all
- the resources accessed by a compute pipeline. It includes zero or more
- `DescriptorSetLayout`s and (optional) push constants.
-
-One can create `DescriptorSetLayout`, `DescriptorSet`, and `ExecutableLayout`
-objects beforehand to avoid incurring overhead during tight computing loops
-and also amortize costs by sharing these objects. However, this isn't totally
-matching Metal's paradigm.
-
-In the Metal framework, the closest concept to `DescriptorSet` would be [argument
-buffer][mtl-argument-buffer]. There is no direct correspondence to
-`DescriptorSetLayout` and `ExecutableLayout`. Rather, the layout is implicitly
-encoded in Metal shaders as MSL structs. The APIs for creating argument buffers
-do not encourage early creation without pipelines: one typically creates them
-for each `MTLFunction`. Besides, unlike Vulkan where different descriptor sets
-can have the same binding number, in Metal even if we have multiple argument
-buffers, the indices for resources are in the same namespace and are typically
-assigned sequentially. That means we need to remap `DescriptorSet`s with a set
-number greater than zero by applying an offset to each of its bindings.
-
-All of this means it's better to defer the creation of the argument buffer
-until the point of compute pipeline creation and dispatch. Therefore, although
-the Metal HAL driver provides the implementation for `DescriptorSet`
-(i.e., `hal::metal::MetalArgumentBuffer`), `DescriptorSetLayout` (i.e.,
-`hal::metal::MetalArgumentBufferLayout`), and `ExecutableLayout` (i.e.,
-`hal::metal::MetalPipelineArgumentBufferLayout`), they are just containers
-holding the information up until the [command buffer
-dispatch](#command-buffer-dispatch) time.
-
-With the above said, the overall idea is still to map one descriptor set to one
-argument buffer. It just means we need to condense and remap the bindings.
-
-### Command buffer dispatch
-
-`MetalCommandBuffer::Dispatch()` performs the following steps with the current
-active `MTLComputeCommandEncoder`:
-
-1. Bind the `MTLComputePipelineState` for the current entry function queried
- from `MetalKernelLibrary`.
-1. For each bound descriptor set at set #`S`:
- 1. Create a [`MTLArgumentEncoder`][mtl-argument-encoder] for encoding an
- associated argument `MTLBuffer`.
- 1. For each bound resource buffer at binding #`B` in this descriptor set,
- encode it to the argument buffer index #`B` with
- `setBuffer::offset::atIndex:` and inform the `MTLComputeCommandEncoder`
- that the dispatch will use this resource with `useResource:usage:`.
- 1. Set the argument `MTLBuffer` to buffer index #`S`.
-1. Dispatch with `dispatchThreadgroups:threadsPerThreadgroup:`.
-
-(TODO: condense and remap bindings)
-
-## Memory Management
-
-### Storage type
-
-Metal provides four [`MTLStorageMode`][mtl-storage-mode] options:
-
-* `MTLStorageModeShared`: The resource is stored in system memory and is
- accessible to both the CPU and the GPU.
-* `MTLStorageModeManaged`: The CPU and GPU may maintain separate copies of the
- resource, and any changes must be explicitly synchronized.
-* `MTLStorageModePrivate`: The resource can be accessed only by the GPU.
-* `MTLStorageMemoryless`: The resource’s contents can be accessed only by the
- GPU and only exist temporarily during a render pass.
-
-Among them, `MTLStorageModeManaged` is only available on macOS.
-
-IREE HAL defines serveral [`MemoryType`][hal-buffer]. They need to map to the
-above storage modes:
-
-* If `kDeviceLocal` but not `kHostVisible`, `MTLStorageModePrivate` is chosen.
-* If `kDeviceLocal` and `kHostVisible`:
- * If macOS, `MTLStorageModeManaged` can be chosen.
- * Otherwise, `MTLStorageModeShared` is chosen.
-* If not `DeviceLocal` but `kDeviceVisible`, `MTLStorageModeShared` is chosen.
-* If not `kDeviceLocal` and not `kDeviceVisible`, `MTLStorageModeShared` is
- chosen. (TODO: We should probably use host buffer here.)
-
-IREE HAL also allows to create buffers with `kHostCoherent` bit. This may still
-be backed by `MTLStorageModeManaged` `MTLBuffer`s in macOS. To respect the
-`kHostCoherent` protocol, the Metal HAL driver will perform necessary
-`InValidate`/`Flush` operations automatically under the hood.
-
-[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/google/iree/tree/main/iree/hal
-[iree-metal]: https://github.com/google/iree/tree/main/iree/hal/metal
-[iree-refptr]: https://github.com/google/iree/blob/main/iree/base/ref_ptr.h
-[hal-allocator]: https://github.com/google/iree/blob/main/iree/hal/allocator.h
-[hal-buffer]: https://github.com/google/iree/blob/main/iree/hal/buffer.h
-[hal-command-queue]: https://github.com/google/iree/blob/main/iree/hal/command_queue.h
-[hal-command-buffer]: https://github.com/google/iree/blob/main/iree/hal/command_buffer.h
-[hal-descriptor-set]: https://github.com/google/iree/blob/main/iree/hal/descriptor_set.h
-[hal-descriptor-set-layout]: https://github.com/google/iree/blob/main/iree/hal/descriptor_set_layout.h
-[hal-executable-layout]: https://github.com/google/iree/blob/main/iree/hal/executable_layout.h
-[hal-device]: https://github.com/google/iree/blob/main/iree/hal/device.h
-[hal-driver]: https://github.com/google/iree/blob/main/iree/hal/driver.h
-[hal-executable]: https://github.com/google/iree/blob/main/iree/hal/executable.h
-[hal-executable-cache]: https://github.com/google/iree/blob/main/iree/hal/executable_cache.h
-[hal-semaphore]: https://github.com/google/iree/blob/main/iree/hal/semaphore.h
-[metal-command-queue]: https://github.com/google/iree/blob/main/iree/hal/metal/metal_command_queue.h
-[metal-command-buffer]: https://github.com/google/iree/blob/main/iree/hal/metal/metal_command_buffer.h
-[metal-device]: https://github.com/google/iree/blob/main/iree/hal/metal/metal_device.h
-[metal-driver]: https://github.com/google/iree/blob/main/iree/hal/metal/metal_driver.h
-[metal-kernel-library]: https://github.com/google/iree/blob/main/iree/hal/metal/metal_kernel_library.h
-[metal-shared-event]: https://github.com/google/iree/blob/main/iree/hal/metal/metal_shared_event.h
-[metal-spirv-target]: https://github.com/google/iree/tree/hal-metal/iree/compiler/Dialect/HAL/Target/MetalSPIRV
-[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
-[mtl-command-buffer]: https://developer.apple.com/documentation/metal/mtlcommandbuffer?language=objc
-[mtl-command-encoder]: https://developer.apple.com/documentation/metal/mtlcommandencoder?language=objc
-[mtl-command-queue]: https://developer.apple.com/documentation/metal/mtlcommandqueue?language=objc
-[mtl-device]: https://developer.apple.com/documentation/metal/mtldevice?language=objc
-[mtl-function]: https://developer.apple.com/documentation/metal/mtlfunction?language=objc
-[mtl-heap]: https://developer.apple.com/documentation/metal/mtlheap?language=objc
-[mtl-library]: https://developer.apple.com/documentation/metal/mtllibrary?language=objc
-[mtl-shared-event]: https://developer.apple.com/documentation/metal/mtlsharedevent?language=objc
-[mtl-storage-mode]: https://developer.apple.com/documentation/metal/mtlstoragemode?language=objc
-[msl-spec]: https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf
-[msl-cl-library]: https://developer.apple.com/documentation/metal/libraries/building_a_library_with_metal_s_command-line_tools?language=objc
-[objc-arc]: https://en.wikipedia.org/wiki/Automatic_Reference_Counting
-[objcxx]: https://en.wikipedia.org/wiki/Objective-C#Objective-C++
-[flatbuffer]: https://google.github.io/flatbuffers/
-[mmap]: https://en.wikipedia.org/wiki/Mmap
-[moltenvk]: https://github.com/KhronosGroup/MoltenVK
-[spirv-codegen]: https://google.github.io/iree/design-docs/codegen-passes
-[spirv-cross]: https://github.com/KhronosGroup/SPIRV-Cross
-[vma]: https://github.com/GPUOpen-LibrariesAndSDKs/VulkanMemoryAllocator
-[vulkan-spirv-target]: https://github.com/google/iree/tree/hal-metal/iree/compiler/Dialect/HAL/Target/VulkanSPIRV
-[vulkan-cmd-dispatch]: https://www.khronos.org/registry/vulkan/specs/1.2-extensions/man/html/vkCmdDispatch.html
diff --git a/iree/hal/drivers/BUILD b/iree/hal/drivers/BUILD
index 4bf772b..8c421d7 100644
--- a/iree/hal/drivers/BUILD
+++ b/iree/hal/drivers/BUILD
@@ -30,8 +30,5 @@
"//iree/hal/dylib/registration",
"//iree/hal/vmla/registration",
"//iree/hal/vulkan/registration",
- ] + select({
- "@bazel_tools//src/conditions:darwin": ["//iree/hal/metal/registration"],
- "//conditions:default": [],
- }),
+ ],
)
diff --git a/iree/hal/drivers/CMakeLists.txt b/iree/hal/drivers/CMakeLists.txt
index 3671c84..305bcb5 100644
--- a/iree/hal/drivers/CMakeLists.txt
+++ b/iree/hal/drivers/CMakeLists.txt
@@ -18,9 +18,6 @@
if(${IREE_HAL_DRIVER_DYLIB})
list(APPEND IREE_HAL_DRIVER_MODULES iree::hal::dylib::registration)
endif()
-if(${IREE_HAL_DRIVER_METAL})
- list(APPEND IREE_HAL_DRIVER_MODULES iree::hal::metal::registration)
-endif()
if(${IREE_HAL_DRIVER_VMLA})
list(APPEND IREE_HAL_DRIVER_MODULES iree::hal::vmla::registration)
endif()
diff --git a/iree/hal/drivers/init.c b/iree/hal/drivers/init.c
index 2b9a2b9..bb36616 100644
--- a/iree/hal/drivers/init.c
+++ b/iree/hal/drivers/init.c
@@ -20,10 +20,6 @@
#include "iree/hal/dylib/registration/driver_module.h"
#endif // IREE_HAL_HAVE_DYLIB_DRIVER_MODULE
-#if defined(IREE_HAL_HAVE_METAL_DRIVER_MODULE)
-#include "iree/hal/metal/registration/driver_module.h"
-#endif // IREE_HAL_HAVE_METAL_DRIVER_MODULE
-
#if defined(IREE_HAL_HAVE_VMLA_DRIVER_MODULE)
#include "iree/hal/vmla/registration/driver_module.h"
#endif // IREE_HAL_HAVE_VMLA_DRIVER_MODULE
@@ -41,11 +37,6 @@
z0, iree_hal_dylib_driver_module_register(registry));
#endif // IREE_HAL_HAVE_DYLIB_DRIVER_MODULE
-#if defined(IREE_HAL_HAVE_METAL_DRIVER_MODULE)
- IREE_RETURN_AND_END_ZONE_IF_ERROR(
- z0, iree_hal_metal_driver_module_register(registry));
-#endif // IREE_HAL_HAVE_METAL_DRIVER_MODULE
-
#if defined(IREE_HAL_HAVE_VMLA_DRIVER_MODULE)
IREE_RETURN_AND_END_ZONE_IF_ERROR(
z0, iree_hal_vmla_driver_module_register(registry));
diff --git a/iree/hal/metal/BUILD.bazel b/iree/hal/metal/BUILD.bazel
deleted file mode 100644
index 388b824..0000000
--- a/iree/hal/metal/BUILD.bazel
+++ /dev/null
@@ -1,70 +0,0 @@
-# Copyright 2020 Google LLC
-#
-# Licensed under the Apache License, Version 2.0 (the "License");
-# you may not use this file except in compliance with the License.
-# You may obtain a copy of the License at
-#
-# https://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing, software
-# distributed under the License is distributed on an "AS IS" BASIS,
-# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-# See the License for the specific language governing permissions and
-# limitations under the License.
-
-package(
- default_visibility = ["//visibility:public"],
- features = ["layering_check"],
- licenses = ["notice"], # Apache 2.0
-)
-
-objc_library(
- name = "metal",
- srcs = [
- "metal_pipeline_argument_buffer.cc",
- ],
- hdrs = [
- "dispatch_time_util.h",
- "metal_buffer.h",
- "metal_capture_manager.h",
- "metal_command_buffer.h",
- "metal_command_queue.h",
- "metal_device.h",
- "metal_direct_allocator.h",
- "metal_driver.h",
- "metal_kernel_library.h",
- "metal_pipeline_argument_buffer.h",
- "metal_pipeline_cache.h",
- "metal_shared_event.h",
- ],
- copts = ["-std=c++14"],
- non_arc_srcs = [
- "metal_buffer.mm",
- "metal_capture_manager.mm",
- "metal_command_buffer.mm",
- "metal_command_queue.mm",
- "metal_device.mm",
- "metal_direct_allocator.mm",
- "metal_driver.mm",
- "metal_kernel_library.mm",
- "metal_pipeline_cache.mm",
- "metal_shared_event.mm",
- ],
- sdk_frameworks = [
- "Foundation",
- "Metal",
- ],
- deps = [
- "//iree/base:arena",
- "//iree/base:file_io",
- "//iree/base:flatcc",
- "//iree/base:logging",
- "//iree/base:status",
- "//iree/base:tracing",
- "//iree/hal",
- "//iree/hal:command_buffer_validation",
- "//iree/schemas:metal_executable_def_c_fbs",
- "@com_google_absl//absl/container:flat_hash_map",
- "@com_google_absl//absl/container:inlined_vector",
- ],
-)
diff --git a/iree/hal/metal/CMakeLists.txt b/iree/hal/metal/CMakeLists.txt
deleted file mode 100644
index 7106f07..0000000
--- a/iree/hal/metal/CMakeLists.txt
+++ /dev/null
@@ -1,66 +0,0 @@
-# Copyright 2020 Google LLC
-#
-# Licensed under the Apache License, Version 2.0 (the "License");
-# you may not use this file except in compliance with the License.
-# You may obtain a copy of the License at
-#
-# https://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing, software
-# distributed under the License is distributed on an "AS IS" BASIS,
-# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-# See the License for the specific language governing permissions and
-# limitations under the License.
-
-if(NOT ${IREE_HAL_DRIVER_METAL})
- return()
-endif()
-
-iree_add_all_subdirs()
-
-iree_cc_library(
- NAME
- metal
- HDRS
- "metal_buffer.h"
- "metal_capture_manager.h"
- "metal_command_buffer.h"
- "metal_command_queue.h"
- "metal_device.h"
- "metal_direct_allocator.h"
- "metal_driver.h"
- "metal_kernel_library.h"
- "metal_pipeline_argument_buffer.h"
- "metal_pipeline_cache.h"
- "metal_shared_event.h"
- SRCS
- "metal_buffer.mm"
- "metal_capture_manager.mm"
- "metal_command_buffer.mm"
- "metal_command_queue.mm"
- "metal_device.mm"
- "metal_direct_allocator.mm"
- "metal_driver.mm"
- "metal_kernel_library.mm"
- "metal_pipeline_argument_buffer.cc"
- "metal_pipeline_cache.mm"
- "metal_shared_event.mm"
- DEPS
- absl::flat_hash_map
- absl::inlined_vector
- absl::memory
- absl::span
- absl::strings
- iree::base::flatcc
- iree::base::file_io
- iree::base::logging
- iree::base::status
- iree::base::time
- iree::base::tracing
- iree::hal::cc
- iree::schemas::metal_executable_def_c_fbs
- LINKOPTS
- "-framework Foundation"
- "-framework Metal"
- PUBLIC
-)
diff --git a/iree/hal/metal/README.md b/iree/hal/metal/README.md
deleted file mode 100644
index d0bbc85..0000000
--- a/iree/hal/metal/README.md
+++ /dev/null
@@ -1,7 +0,0 @@
-# Metal HAL Driver
-
-**TODO(antiagainst)**: move the docs here - having them separate is suboptimal.
-
-This directory contains the source code for the Metal HAL driver. See the
-[design doc](https://google.github.io/iree/design-docs/metal-hal-driver) for
-more details.
diff --git a/iree/hal/metal/dispatch_time_util.h b/iree/hal/metal/dispatch_time_util.h
deleted file mode 100644
index 6023d38..0000000
--- a/iree/hal/metal/dispatch_time_util.h
+++ /dev/null
@@ -1,44 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#ifndef IREE_HAL_METAL_APPLE_TIME_UTIL_H_
-#define IREE_HAL_METAL_APPLE_TIME_UTIL_H_
-
-#include <dispatch/dispatch.h>
-
-#include "iree/base/time.h"
-
-namespace iree {
-namespace hal {
-namespace metal {
-
-// Converts a relative iree::Duration against the currrent time to the
-// corresponding dispatch_time_t value.
-static inline dispatch_time_t DurationToDispatchTime(Duration duration_ns) {
- if (duration_ns == InfiniteDuration()) return DISPATCH_TIME_FOREVER;
- if (duration_ns == ZeroDuration()) return DISPATCH_TIME_NOW;
- return dispatch_time(DISPATCH_TIME_NOW, static_cast<uint64_t>(duration_ns));
-}
-
-// Converts an absolute iree::Time time to the corresponding dispatch_time_t
-// value.
-static inline dispatch_time_t DeadlineToDispatchTime(Time deadline_ns) {
- return DurationToDispatchTime(DeadlineToRelativeTimeoutNanos(deadline_ns));
-}
-
-} // namespace metal
-} // namespace hal
-} // namespace iree
-
-#endif // IREE_HAL_METAL_APPLE_TIME_UTIL_H_
diff --git a/iree/hal/metal/metal_buffer.h b/iree/hal/metal/metal_buffer.h
deleted file mode 100644
index 4238d51..0000000
--- a/iree/hal/metal/metal_buffer.h
+++ /dev/null
@@ -1,101 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#ifndef IREE_HAL_METAL_METAL_BUFFER_H_
-#define IREE_HAL_METAL_METAL_BUFFER_H_
-
-#import <Metal/Metal.h>
-
-#include "iree/hal/api.h"
-
-id<MTLBuffer> iree_hal_metal_buffer_handle(iree_hal_buffer_t* base_buffer);
-
-namespace iree {
-namespace hal {
-namespace metal {
-
-class MetalDirectAllocator;
-
-// A buffer implementation for Metal that directly wraps a MTLBuffer.
-class MetalBuffer final : public Buffer {
- public:
- // Creates a MetalBuffer instance with retaining the given id<MTLBuffer>.
- static StatusOr<ref_ptr<MetalBuffer>> Create(
- MetalDirectAllocator* allocator, iree_hal_memory_type_t memory_type,
- iree_hal_memory_access_t allowed_access, iree_hal_buffer_usage_t usage,
- iree_device_size_t allocation_size, iree_device_size_t byte_offset,
- iree_device_size_t byte_length, id<MTLBuffer> buffer,
- id<MTLCommandQueue> transfer_queue);
-
- ~MetalBuffer() override;
-
- id<MTLBuffer> handle() const { return metal_handle_; }
-
- private:
- // Creates a MetalBuffer instance without retaining the given id<MTLBuffer>.
- MetalBuffer(MetalDirectAllocator* allocator,
- iree_hal_memory_type_t memory_type,
- iree_hal_memory_access_t allowed_access,
- iree_hal_buffer_usage_t usage, iree_device_size_t allocation_size,
- iree_device_size_t byte_offset, iree_device_size_t byte_length,
- id<MTLBuffer> buffer, id<MTLCommandQueue> transfer_queue);
-
- Status FillImpl(iree_device_size_t byte_offset,
- iree_device_size_t byte_length, const void* pattern,
- iree_device_size_t pattern_length) override;
- Status ReadDataImpl(iree_device_size_t source_offset, void* data,
- iree_device_size_t data_length) override;
- Status WriteDataImpl(iree_device_size_t target_offset, const void* data,
- iree_device_size_t data_length) override;
- Status CopyDataImpl(iree_device_size_t target_offset, Buffer* source_buffer,
- iree_device_size_t source_offset,
- iree_device_size_t data_length) override;
-
- Status MapMemoryImpl(MappingMode mapping_mode,
- iree_hal_memory_access_t memory_access,
- iree_device_size_t local_byte_offset,
- iree_device_size_t local_byte_length,
- void** out_data) override;
- Status UnmapMemoryImpl(iree_device_size_t local_byte_offset,
- iree_device_size_t local_byte_length,
- void* data) override;
- Status InvalidateMappedMemoryImpl(
- iree_device_size_t local_byte_offset,
- iree_device_size_t local_byte_length) override;
- Status FlushMappedMemoryImpl(iree_device_size_t local_byte_offset,
- iree_device_size_t local_byte_length) override;
-
- // Returns true if we need to automatically invaliate/flush CPU caches to keep
- // memory hierarchy consistent.
- //
- // Note: this is needed when the buffer is requested with
- // IREE_HAL_MEMORY_TYPE_HOST_COHERENT bit but under the hood we are using
- // memory types that does not have that property natively, e.g.,
- // MTLStorageModeManaged. Under such circumstances, we need to perform the
- // invalidate/flush operation "automatically" for users.
- bool requires_autosync() const;
-
- // We need to hold an reference to the queue so that we can encode
- // synchronizeResource commands for synchronizing the buffer with
- // MTLResourceStorageModeManaged.
- id<MTLCommandQueue> metal_transfer_queue_;
-
- id<MTLBuffer> metal_handle_;
-};
-
-} // namespace metal
-} // namespace hal
-} // namespace iree
-
-#endif // IREE_HAL_METAL_METAL_BUFFER_H_
diff --git a/iree/hal/metal/metal_buffer.mm b/iree/hal/metal/metal_buffer.mm
deleted file mode 100644
index a6cb965..0000000
--- a/iree/hal/metal/metal_buffer.mm
+++ /dev/null
@@ -1,194 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#include "iree/hal/metal/metal_buffer.h"
-
-#include "iree/base/status.h"
-#include "iree/base/tracing.h"
-#include "iree/hal/metal/metal_direct_allocator.h"
-
-id<MTLBuffer> iree_hal_metal_buffer_handle(iree_hal_buffer_t* base_buffer);
-
-namespace iree {
-namespace hal {
-namespace metal {
-
-// static
-StatusOr<ref_ptr<MetalBuffer>> MetalBuffer::Create(
- MetalDirectAllocator* allocator, iree_hal_memory_type_t memory_type,
- iree_hal_memory_access_t allowed_access, iree_hal_buffer_usage_t usage, iree_device_size_t allocation_size,
- iree_device_size_t byte_offset, iree_device_size_t byte_length, id<MTLBuffer> buffer,
- id<MTLCommandQueue> transfer_queue) {
- IREE_TRACE_SCOPE0("MetalBuffer::Create");
- return assign_ref(new MetalBuffer(allocator, memory_type, allowed_access, usage, allocation_size,
- byte_offset, byte_length, buffer, transfer_queue));
-}
-
-MetalBuffer::MetalBuffer(MetalDirectAllocator* allocator, iree_hal_memory_type_t memory_type,
- iree_hal_memory_access_t allowed_access, iree_hal_buffer_usage_t usage,
- iree_device_size_t allocation_size, iree_device_size_t byte_offset,
- iree_device_size_t byte_length, id<MTLBuffer> buffer,
- id<MTLCommandQueue> transfer_queue)
- : Buffer(allocator, memory_type, allowed_access, usage, allocation_size, byte_offset,
- byte_length),
- metal_transfer_queue_([transfer_queue retain]),
- metal_handle_(buffer) {}
-
-MetalBuffer::~MetalBuffer() {
- IREE_TRACE_SCOPE0("MetalBuffer::dtor");
- [metal_handle_ release];
- [metal_transfer_queue_ release];
-}
-
-Status MetalBuffer::FillImpl(iree_device_size_t byte_offset, iree_device_size_t byte_length,
- const void* pattern, iree_device_size_t pattern_length) {
- IREE_ASSIGN_OR_RETURN(auto mapping,
- MapMemory<uint8_t>(IREE_HAL_MEMORY_ACCESS_DISCARD_WRITE, byte_offset, byte_length));
- void* data_ptr = static_cast<void*>(mapping.mutable_data());
- switch (pattern_length) {
- case 1: {
- uint8_t* data = static_cast<uint8_t*>(data_ptr);
- uint8_t value_bits = *static_cast<const uint8_t*>(pattern);
- std::fill_n(data, byte_length, value_bits);
- break;
- }
- case 2: {
- uint16_t* data = static_cast<uint16_t*>(data_ptr);
- uint16_t value_bits = *static_cast<const uint16_t*>(pattern);
- std::fill_n(data, byte_length / sizeof(uint16_t), value_bits);
- break;
- }
- case 4: {
- uint32_t* data = static_cast<uint32_t*>(data_ptr);
- uint32_t value_bits = *static_cast<const uint32_t*>(pattern);
- std::fill_n(data, byte_length / sizeof(uint32_t), value_bits);
- break;
- }
- default:
- return InvalidArgumentErrorBuilder(IREE_LOC)
- << "Unsupported scalar data size: " << pattern_length;
- }
- return OkStatus();
-}
-
-Status MetalBuffer::ReadDataImpl(iree_device_size_t source_offset, void* data,
- iree_device_size_t data_length) {
- IREE_ASSIGN_OR_RETURN(auto mapping,
- MapMemory<uint8_t>(IREE_HAL_MEMORY_ACCESS_READ, source_offset, data_length));
- std::memcpy(data, mapping.data(), mapping.byte_length());
- return OkStatus();
-}
-
-Status MetalBuffer::WriteDataImpl(iree_device_size_t target_offset, const void* data,
- iree_device_size_t data_length) {
- IREE_ASSIGN_OR_RETURN(
- auto mapping, MapMemory<uint8_t>(IREE_HAL_MEMORY_ACCESS_DISCARD_WRITE, target_offset, data_length));
- std::memcpy(mapping.mutable_data(), data, mapping.byte_length());
- return OkStatus();
-}
-
-Status MetalBuffer::CopyDataImpl(iree_device_size_t target_offset, Buffer* source_buffer,
- iree_device_size_t source_offset, iree_device_size_t data_length) {
- // This is pretty terrible. Let's not do this.
- // TODO(benvanik): a way for allocators to indicate transfer compat.
- IREE_ASSIGN_OR_RETURN(auto source_mapping, source_buffer->MapMemory<uint8_t>(
- IREE_HAL_MEMORY_ACCESS_READ, source_offset, data_length));
- IREE_CHECK_EQ(data_length, source_mapping.size());
- IREE_ASSIGN_OR_RETURN(auto target_mapping, MapMemory<uint8_t>(IREE_HAL_MEMORY_ACCESS_DISCARD_WRITE,
- target_offset, data_length));
- IREE_CHECK_EQ(data_length, target_mapping.size());
- std::memcpy(target_mapping.mutable_data(), source_mapping.data(), data_length);
- return OkStatus();
-}
-
-Status MetalBuffer::MapMemoryImpl(MappingMode mapping_mode, iree_hal_memory_access_t memory_access,
- iree_device_size_t local_byte_offset, iree_device_size_t local_byte_length,
- void** out_data) {
- uint8_t* data_ptr = reinterpret_cast<uint8_t*>([metal_handle_ contents]);
- *out_data = data_ptr + local_byte_offset;
-
- // If we mapped for discard scribble over the bytes. This is not a mandated
- // behavior but it will make debugging issues easier. Alternatively for
- // heap buffers we could reallocate them such that ASAN yells, but that
- // would only work if the entire buffer was discarded.
-#ifndef NDEBUG
- if (iree_any_bit_set(memory_access, IREE_HAL_MEMORY_ACCESS_DISCARD)) {
- std::memset(data_ptr + local_byte_offset, 0xCD, local_byte_length);
- }
-#endif // !NDEBUG
-
- if (requires_autosync()) {
- IREE_RETURN_IF_ERROR(InvalidateMappedMemoryImpl(local_byte_offset, local_byte_length));
- }
-
- return OkStatus();
-}
-
-Status MetalBuffer::UnmapMemoryImpl(iree_device_size_t local_byte_offset,
- iree_device_size_t local_byte_length, void* data) {
- if (requires_autosync()) {
- IREE_RETURN_IF_ERROR(FlushMappedMemoryImpl(local_byte_offset, local_byte_length));
- }
-
- return OkStatus();
-}
-
-Status MetalBuffer::InvalidateMappedMemoryImpl(iree_device_size_t local_byte_offset,
- iree_device_size_t local_byte_length) {
-#ifdef IREE_PLATFORM_MACOS
- // The following is only necessary for MTLStorageManaged.
- if (metal_handle_.storageMode == MTLStorageModeManaged) {
- @autoreleasepool {
- id<MTLCommandBuffer> command_buffer =
- [metal_transfer_queue_ commandBufferWithUnretainedReferences];
-
- id<MTLBlitCommandEncoder> blit_encoder = [command_buffer blitCommandEncoder];
- [blit_encoder synchronizeResource:metal_handle_];
- [blit_encoder endEncoding];
-
- [command_buffer commit];
- [command_buffer waitUntilCompleted];
- }
- }
-#endif
-
- return OkStatus();
-}
-
-Status MetalBuffer::FlushMappedMemoryImpl(iree_device_size_t local_byte_offset,
- iree_device_size_t local_byte_length) {
-#ifdef IREE_PLATFORM_MACOS
- // The following is only necessary for MTLStorageManaged.
- if (metal_handle_.storageMode == MTLStorageModeManaged) {
- [metal_handle_ didModifyRange:NSMakeRange(local_byte_offset, local_byte_length)];
- }
-#endif
-
- return OkStatus();
-}
-
-bool MetalBuffer::requires_autosync() const {
- // We only need to perform "automatic" resource synchronization if it's MTLStorageModeManaged,
- // which is only available on macOS.
-#ifdef IREE_PLATFORM_MACOS
- return iree_all_bits_set(memory_type(), IREE_HAL_MEMORY_TYPE_HOST_COHERENT) &&
- metal_handle_.storageMode == MTLStorageModeManaged;
-#else
- return false;
-#endif
-}
-
-} // namespace metal
-} // namespace hal
-} // namespace iree
diff --git a/iree/hal/metal/metal_capture_manager.h b/iree/hal/metal/metal_capture_manager.h
deleted file mode 100644
index 22d17dd..0000000
--- a/iree/hal/metal/metal_capture_manager.h
+++ /dev/null
@@ -1,64 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#ifndef IREE_HAL_METAL_METAL_CAPTURE_MANAGER_H_
-#define IREE_HAL_METAL_METAL_CAPTURE_MANAGER_H_
-
-#include <memory>
-
-#import <Metal/Metal.h>
-
-#include "iree/base/status.h"
-
-namespace iree {
-namespace hal {
-namespace metal {
-
-// A DebugCaptureManager implementation for Metal that directly wraps a
-// MTLCaptureManager.
-class MetalCaptureManager final : public DebugCaptureManager {
- public:
- // Creates a capture manager that captures Metal commands to the given |capture_file| if not
- // empty. Capture to Xcode otherwise.
- static StatusOr<std::unique_ptr<MetalCaptureManager>> Create(const std::string& capture_file);
- ~MetalCaptureManager() override;
-
- Status Connect() override;
-
- void Disconnect() override;
-
- bool is_connected() const override;
-
- void SetCaptureObject(id object);
-
- void StartCapture() override;
-
- void StopCapture() override;
-
- bool is_capturing() const override;
-
- private:
- explicit MetalCaptureManager(NSURL* capture_file);
-
- MTLCaptureManager* metal_handle_ = nil;
- // The path for storing the .gputrace file. Empty means capturing to Xcode.
- NSURL* capture_file_ = nil;
- id capture_object_ = nil;
-};
-
-} // namespace metal
-} // namespace hal
-} // namespace iree
-
-#endif // IREE_HAL_METAL_METAL_CAPTURE_MANAGER_H_
diff --git a/iree/hal/metal/metal_capture_manager.mm b/iree/hal/metal/metal_capture_manager.mm
deleted file mode 100644
index 4437951..0000000
--- a/iree/hal/metal/metal_capture_manager.mm
+++ /dev/null
@@ -1,128 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#include "iree/hal/metal/metal_capture_manager.h"
-
-#include <string>
-
-#include "absl/memory/memory.h"
-#include "iree/base/file_io.h"
-#include "iree/base/logging.h"
-#include "iree/base/tracing.h"
-
-namespace iree {
-namespace hal {
-namespace metal {
-
-// static
-StatusOr<std::unique_ptr<MetalCaptureManager>> MetalCaptureManager::Create(
- const std::string& capture_file) {
- IREE_TRACE_SCOPE0("MetalCaptureManager::Create");
- @autoreleasepool {
- NSURL* capture_url = nil;
- if (!capture_file.empty()) {
- NSString* ns_string = [NSString stringWithCString:capture_file.c_str()
- encoding:[NSString defaultCStringEncoding]];
- NSString* capture_path = ns_string.stringByStandardizingPath;
- capture_url = [[NSURL fileURLWithPath:capture_path isDirectory:false] retain];
- }
- return absl::WrapUnique(new MetalCaptureManager(capture_url));
- }
-}
-
-MetalCaptureManager::MetalCaptureManager(NSURL* capture_file) : capture_file_(capture_file) {}
-
-MetalCaptureManager::~MetalCaptureManager() {
- IREE_TRACE_SCOPE0("MetalCaptureManager::dtor");
- Disconnect();
- if (capture_file_) [capture_file_ release];
-}
-
-Status MetalCaptureManager::Connect() {
- IREE_TRACE_SCOPE0("MetalCaptureManager::Connect");
-
- if (metal_handle_) return OkStatus();
-
- @autoreleasepool {
- metal_handle_ = [[MTLCaptureManager sharedCaptureManager] retain];
-
- if (capture_file_ &&
- [metal_handle_ supportsDestination:MTLCaptureDestinationGPUTraceDocument]) {
- IREE_LOG(INFO) << "Connected to shared Metal capture manager; writing capture to "
- << std::string([capture_file_.absoluteString UTF8String]);
- } else {
- IREE_LOG(INFO) << "Connected to shared Metal capture manager; capturing to Xcode";
- }
- }
-
- return OkStatus();
-}
-
-void MetalCaptureManager::Disconnect() {
- IREE_TRACE_SCOPE0("MetalCaptureManager::Disconnect");
-
- if (!metal_handle_) return;
-
- if (is_capturing()) StopCapture();
-
- [metal_handle_ release];
- metal_handle_ = nil;
-}
-
-bool MetalCaptureManager::is_connected() const { return metal_handle_ != nil; }
-
-void MetalCaptureManager::SetCaptureObject(id object) { capture_object_ = object; }
-
-void MetalCaptureManager::StartCapture() {
- IREE_TRACE_SCOPE0("MetalCaptureManager::StartCapture");
-
- IREE_CHECK(is_connected()) << "Can't start capture when not connected";
- IREE_CHECK(!is_capturing()) << "Capture is already started";
- IREE_CHECK(capture_object_) << "Must set capture object before starting";
-
- IREE_LOG(INFO) << "Starting Metal capture";
- @autoreleasepool {
- MTLCaptureDescriptor* capture_descriptor = [[[MTLCaptureDescriptor alloc] init] autorelease];
- capture_descriptor.captureObject = capture_object_;
- if (capture_file_) {
- capture_descriptor.destination = MTLCaptureDestinationGPUTraceDocument;
- capture_descriptor.outputURL = capture_file_;
- } else {
- capture_descriptor.destination = MTLCaptureDestinationDeveloperTools;
- }
-
- NSError* error;
- if (![metal_handle_ startCaptureWithDescriptor:capture_descriptor error:&error]) {
- NSLog(@"Failed to start capture, error %@", error);
- }
- }
-}
-
-void MetalCaptureManager::StopCapture() {
- IREE_TRACE_SCOPE0("MetalCaptureManager::StopCapture");
-
- IREE_CHECK(is_capturing()) << "Can't stop capture when not capturing";
-
- IREE_LOG(INFO) << "Ending Metal capture";
- [metal_handle_ stopCapture];
-}
-
-bool MetalCaptureManager::is_capturing() const {
- if (!is_connected()) return false;
- return metal_handle_.isCapturing;
-}
-
-} // namespace metal
-} // namespace hal
-} // namespace iree
diff --git a/iree/hal/metal/metal_command_buffer.h b/iree/hal/metal/metal_command_buffer.h
deleted file mode 100644
index d9b4321..0000000
--- a/iree/hal/metal/metal_command_buffer.h
+++ /dev/null
@@ -1,148 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#ifndef IREE_HAL_METAL_METAL_COMMAND_BUFFER_H_
-#define IREE_HAL_METAL_METAL_COMMAND_BUFFER_H_
-
-#import <Metal/Metal.h>
-
-#include "absl/container/flat_hash_map.h"
-#include "absl/container/inlined_vector.h"
-#include "iree/hal/metal/metal_buffer.h"
-
-namespace iree {
-namespace hal {
-namespace metal {
-
-// A command buffer implementation for Metal that directly wraps a
-// MTLCommandBuffer.
-//
-// Objects of this class are not expected to be accessed by multiple threads.
-class MetalCommandBuffer final : public CommandBuffer {
- public:
- static StatusOr<ref_ptr<CommandBuffer>> Create(
- iree_hal_command_buffer_mode_t mode,
- iree_hal_command_category_t command_categories,
- id<MTLCommandBuffer> command_buffer);
- ~MetalCommandBuffer() override;
-
- id<MTLCommandBuffer> handle() const { return metal_handle_; }
-
- Status Begin() override;
- Status End() override;
-
- Status ExecutionBarrier(
- iree_hal_execution_stage_t source_stage_mask,
- iree_hal_execution_stage_t target_stage_mask,
- absl::Span<const iree_hal_memory_barrier_t> memory_barriers,
- absl::Span<const iree_hal_buffer_barrier_t> buffer_barriers) override;
-
- Status SignalEvent(iree_hal_event_t* event,
- iree_hal_execution_stage_t source_stage_mask) override;
- Status ResetEvent(iree_hal_event_t* event,
- iree_hal_execution_stage_t source_stage_mask) override;
- Status WaitEvents(
- absl::Span<iree_hal_event_t*> events,
- iree_hal_execution_stage_t source_stage_mask,
- iree_hal_execution_stage_t target_stage_mask,
- absl::Span<const iree_hal_memory_barrier_t> memory_barriers,
- absl::Span<const iree_hal_buffer_barrier_t> buffer_barriers) override;
-
- Status FillBuffer(iree_hal_buffer_t* target_buffer,
- iree_device_size_t target_offset, iree_device_size_t length,
- const void* pattern, size_t pattern_length) override;
- Status DiscardBuffer(iree_hal_buffer_t* buffer) override;
- Status UpdateBuffer(const void* source_buffer,
- iree_device_size_t source_offset,
- iree_hal_buffer_t* target_buffer,
- iree_device_size_t target_offset,
- iree_device_size_t length) override;
- Status CopyBuffer(iree_hal_buffer_t* source_buffer,
- iree_device_size_t source_offset,
- iree_hal_buffer_t* target_buffer,
- iree_device_size_t target_offset,
- iree_device_size_t length) override;
-
- Status PushConstants(iree_hal_executable_layout_t* executable_layout,
- size_t offset,
- absl::Span<const uint32_t> values) override;
-
- Status PushDescriptorSet(
- iree_hal_executable_layout_t* executable_layout, uint32_t set,
- absl::Span<const iree_hal_descriptor_set_binding_t> bindings) override;
- Status BindDescriptorSet(
- iree_hal_executable_layout_t* executable_layout, uint32_t set,
- iree_hal_descriptor_set_t* descriptor_set,
- absl::Span<const iree_device_size_t> dynamic_offsets) override;
-
- Status Dispatch(iree_hal_executable_t* executable, int32_t entry_point,
- std::array<uint32_t, 3> workgroups) override;
- Status DispatchIndirect(iree_hal_executable_t* executable,
- int32_t entry_point,
- iree_hal_buffer_t* workgroups_buffer,
- iree_device_size_t workgroups_offset) override;
-
- private:
- // A struct containing all resources states of the current pipeline.
- struct PipelineStateObject {
- struct PushState {
- absl::InlinedVector<iree_hal_descriptor_set_binding_t, 8>
- resource_bindings;
- };
- // Map from set number to push descriptor states
- absl::flat_hash_map<int32_t, PushState> push_states;
-
- struct BindState {
- DescriptorSet* descriptor_set;
- };
- // Map from set number to bind descriptor states
- absl::flat_hash_map<int32_t, BindState> bind_states;
-
- struct ConstantState {
- absl::InlinedVector<uint32_t, 16> values;
- };
- // Map from set number to push constant states
- absl::flat_hash_map<uint32_t, ConstantState> constant_states;
- };
-
- MetalCommandBuffer(iree_hal_command_buffer_mode_t mode,
- iree_hal_command_category_t command_categories,
- id<MTLCommandBuffer> command_buffer);
-
- // Gets or begins an active MTLBlitCommandEncoder. This also ends all previous
- // encoded compute commands if any.
- id<MTLBlitCommandEncoder> GetOrBeginBlitEncoder();
- void EndBlitEncoder();
-
- // Gets or begins a new MTLComputeCommandEncoder. This also ends all previous
- // encoded blit commands if any.
- id<MTLComputeCommandEncoder> GetOrBeginComputeEncoder();
- void EndComputeEncoder();
-
- private:
- bool is_recording_ = false;
- id<MTLCommandBuffer> metal_handle_;
-
- id<MTLComputeCommandEncoder> current_compute_encoder_ = nil;
- id<MTLBlitCommandEncoder> current_blit_encoder_ = nil;
-
- absl::flat_hash_map<iree_hal_executable_layout_t*, PipelineStateObject>
- pipeline_state_objects_;
-};
-
-} // namespace metal
-} // namespace hal
-} // namespace iree
-
-#endif // IREE_HAL_METAL_METAL_COMMAND_BUFFER_H_
diff --git a/iree/hal/metal/metal_command_buffer.mm b/iree/hal/metal/metal_command_buffer.mm
deleted file mode 100644
index 46607f6..0000000
--- a/iree/hal/metal/metal_command_buffer.mm
+++ /dev/null
@@ -1,380 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#include "iree/hal/metal/metal_command_buffer.h"
-
-#include "iree/base/logging.h"
-#include "iree/base/status.h"
-#include "iree/base/tracing.h"
-#include "iree/hal/metal/metal_kernel_library.h"
-#include "iree/hal/metal/metal_pipeline_argument_buffer.h"
-
-namespace iree {
-namespace hal {
-namespace metal {
-
-namespace {
-
-MTLResourceUsage ConvertResourceUsage(iree_hal_memory_access_t memory_access) {
- MTLResourceUsage usage = 0;
- if (iree_all_bits_set(memory_access, IREE_HAL_MEMORY_ACCESS_READ)) usage |= MTLResourceUsageRead;
- if (iree_all_bits_set(memory_access, IREE_HAL_MEMORY_ACCESS_WRITE)) usage |= MTLResourceUsageWrite;
- return usage;
-}
-
-} // namespace
-
-// static
-StatusOr<ref_ptr<CommandBuffer>> MetalCommandBuffer::Create(
- iree_hal_command_buffer_mode_t mode, iree_hal_command_category_t command_categories,
- id<MTLCommandBuffer> command_buffer) {
- return assign_ref(new MetalCommandBuffer(mode, command_categories, command_buffer));
-}
-
-MetalCommandBuffer::MetalCommandBuffer(iree_hal_command_buffer_mode_t mode,
- iree_hal_command_category_t command_categories,
- id<MTLCommandBuffer> command_buffer)
- : CommandBuffer(mode, command_categories), metal_handle_([command_buffer retain]) {
- metal_handle_.label = @"IREE MetalCommandBuffer";
-}
-
-MetalCommandBuffer::~MetalCommandBuffer() {
- IREE_TRACE_SCOPE0("MetalCommandBuffer::dtor");
- [metal_handle_ release];
-}
-
-id<MTLBlitCommandEncoder> MetalCommandBuffer::GetOrBeginBlitEncoder() {
- IREE_TRACE_SCOPE0("MetalCommandBuffer::GetOrBeginBlitEncoder");
-
- if (current_compute_encoder_) EndComputeEncoder();
-
- @autoreleasepool {
- if (!current_blit_encoder_) {
- current_blit_encoder_ = [[metal_handle_ blitCommandEncoder] retain];
- }
- }
-
- return current_blit_encoder_;
-}
-
-void MetalCommandBuffer::EndBlitEncoder() {
- IREE_TRACE_SCOPE0("MetalCommandBuffer::EndBlitEncoder");
- if (current_blit_encoder_) {
- [current_blit_encoder_ endEncoding];
- [current_blit_encoder_ release];
- current_blit_encoder_ = nil;
- }
-}
-
-id<MTLComputeCommandEncoder> MetalCommandBuffer::GetOrBeginComputeEncoder() {
- IREE_TRACE_SCOPE0("MetalCommandBuffer::GetOrBeginComputeEncoder");
-
- if (current_blit_encoder_) EndBlitEncoder();
-
- @autoreleasepool {
- if (!current_compute_encoder_) {
- current_compute_encoder_ = [[metal_handle_ computeCommandEncoder] retain];
- }
- }
-
- return current_compute_encoder_;
-}
-
-void MetalCommandBuffer::EndComputeEncoder() {
- IREE_TRACE_SCOPE0("MetalCommandBuffer::EndComputeEncoder");
- if (current_compute_encoder_) {
- [current_compute_encoder_ endEncoding];
- [current_compute_encoder_ release];
- current_compute_encoder_ = nil;
- }
-}
-
-Status MetalCommandBuffer::Begin() {
- IREE_TRACE_SCOPE0("MetalCommandBuffer::Begin");
- is_recording_ = true;
- return OkStatus();
-}
-
-Status MetalCommandBuffer::End() {
- IREE_TRACE_SCOPE0("MetalCommandBuffer::End");
- EndBlitEncoder();
- EndComputeEncoder();
- is_recording_ = false;
- return OkStatus();
-}
-
-Status MetalCommandBuffer::ExecutionBarrier(iree_hal_execution_stage_t source_stage_mask,
- iree_hal_execution_stage_t target_stage_mask,
- absl::Span<const iree_hal_memory_barrier_t> memory_barriers,
- absl::Span<const iree_hal_buffer_barrier_t> buffer_barriers) {
- IREE_TRACE_SCOPE0("MetalCommandBuffer::ExecutionBarrier");
-
- if (iree_all_bits_set(source_stage_mask, IREE_HAL_EXECUTION_STAGE_HOST) ||
- iree_all_bits_set(target_stage_mask, IREE_HAL_EXECUTION_STAGE_HOST)) {
- return UnimplementedErrorBuilder(IREE_LOC)
- << "MetalCommandBuffer::ExecutionBarrier with host bit set";
- }
-
- // If there is a memory barrier specified, we have to place a catch-all barrier for all buffers.
- // Metal does not provide a more fine-grained control here. But we do have the option to specify a
- // list of buffers to synchronize if only buffer barriers are specified.
- if (!memory_barriers.empty()) {
- [GetOrBeginComputeEncoder() memoryBarrierWithScope:MTLBarrierScopeBuffers];
- } else if (!buffer_barriers.empty()) {
- std::vector<id<MTLResource>> buffers;
- buffers.reserve(buffer_barriers.size());
- for (const auto& barrier : buffer_barriers) {
- buffers.push_back(static_cast<MetalBuffer*>(barrier.buffer)->handle());
- }
- [GetOrBeginComputeEncoder() memoryBarrierWithResources:buffers.data() count:buffers.size()];
- }
-
- return OkStatus();
-}
-
-Status MetalCommandBuffer::SignalEvent(iree_hal_event_t* event, iree_hal_execution_stage_t source_stage_mask) {
- IREE_TRACE_SCOPE0("MetalCommandBuffer::SignalEvent");
- return UnimplementedErrorBuilder(IREE_LOC) << "MetalCommandBuffer::SignalEvent";
-}
-
-Status MetalCommandBuffer::ResetEvent(iree_hal_event_t* event, iree_hal_execution_stage_t source_stage_mask) {
- IREE_TRACE_SCOPE0("MetalCommandBuffer::ResetEvent");
- return UnimplementedErrorBuilder(IREE_LOC) << "MetalCommandBuffer::ResetEvent";
-}
-
-Status MetalCommandBuffer::WaitEvents(absl::Span<iree_hal_event_t*> events,
- iree_hal_execution_stage_t source_stage_mask,
- iree_hal_execution_stage_t target_stage_mask,
- absl::Span<const iree_hal_memory_barrier_t> memory_barriers,
- absl::Span<const iree_hal_buffer_barrier_t> buffer_barriers) {
- IREE_TRACE_SCOPE0("MetalCommandBuffer::WaitEvents");
- return UnimplementedErrorBuilder(IREE_LOC) << "MetalCommandBuffer::WaitEvents";
-}
-
-Status MetalCommandBuffer::FillBuffer(iree_hal_buffer_t* target_buffer, iree_device_size_t target_offset,
- iree_device_size_t length, const void* pattern,
- size_t pattern_length) {
- IREE_TRACE_SCOPE0("MetalCommandBuffer::FillBuffer");
- id<MTLBuffer> target_device_buffer = iree_hal_metal_buffer_handle(iree_hal_buffer_allocated_buffer(target_buffer));
-
- target_offset += iree_hal_buffer_byte_offset(target_buffer);
-
- // Per the spec for fillBuffer:range:value: "The alignment and length of the range must both be a
- // multiple of 4 bytes in macOS, and 1 byte in iOS and tvOS." Although iOS/tvOS is more relaxed on
- // this front, we still require 4-byte alignment for uniformity across IREE.
- if (target_offset % 4 != 0) {
- return UnimplementedErrorBuilder(IREE_LOC)
- << "MetalCommandBuffer::FillBuffer with offset that is not a multiple of 4 bytes";
- }
-
- // Note that fillBuffer:range:value: only accepts a single byte as the pattern but FillBuffer
- // can accept 1/2/4 bytes. If the pattern itself contains repeated bytes, we can call into
- // fillBuffer:range:value:. Otherwise we may need to find another way. Just implement the case
- // where we have a single byte to fill for now.
- if (pattern_length != 1) {
- return UnimplementedErrorBuilder(IREE_LOC)
- << "MetalCommandBuffer::FillBuffer with non-1-byte pattern";
- }
- uint8_t byte_pattern = *reinterpret_cast<const uint8_t*>(pattern);
-
- [GetOrBeginBlitEncoder() fillBuffer:target_device_buffer->handle()
- range:NSMakeRange(target_offset, length)
- value:byte_pattern];
-
- return OkStatus();
-}
-
-Status MetalCommandBuffer::DiscardBuffer(iree_hal_buffer_t* buffer) {
- IREE_TRACE_SCOPE0("MetalCommandBuffer::DiscardBuffer");
- // This is a hint. Nothing to do for Metal.
- return OkStatus();
-}
-
-Status MetalCommandBuffer::UpdateBuffer(const void* source_buffer, iree_device_size_t source_offset,
- iree_hal_buffer_t* target_buffer, iree_device_size_t target_offset,
- iree_device_size_t length) {
- IREE_TRACE_SCOPE0("MetalCommandBuffer::UpdateBuffer");
- return UnimplementedErrorBuilder(IREE_LOC) << "MetalCommandBuffer::UpdateBuffer";
-}
-
-Status MetalCommandBuffer::CopyBuffer(iree_hal_buffer_t* source_buffer, iree_device_size_t source_offset,
- iree_hal_buffer_t* target_buffer, iree_device_size_t target_offset,
- iree_device_size_t length) {
- IREE_TRACE_SCOPE0("MetalCommandBuffer::CopyBuffer");
-
- id<MTLBuffer> source_device_buffer = iree_hal_metal_buffer_handle(iree_hal_buffer_allocated_buffer(source_buffer));
- id<MTLBuffer> target_device_buffer = iree_hal_metal_buffer_handle(iree_hal_buffer_allocated_buffer(target_buffer));
-
- source_offset += iree_hal_buffer_byte_offset(source_buffer);
- target_offset += iree_hal_buffer_byte_offset(target_buffer);
-
- // Per the spec for copyFromBuffer:sourceOffset:toBuffer:destinationOffset:size, the source/target
- // offset must be a multiple of 4 bytes in macOS, and 1 byte in iOS and tvOS. Although iOS/tvOS
- // is more relaxed on this front, we still require 4-byte alignment for uniformity across IREE.
- if (source_offset % 4 != 0 || target_offset % 4 != 0) {
- return UnimplementedErrorBuilder(IREE_LOC)
- << "MetalCommandBuffer::CopyBuffer with offset that is not a multiple of 4 bytes";
- }
-
- [GetOrBeginBlitEncoder() copyFromBuffer:source_device_buffer->handle()
- sourceOffset:source_offset
- toBuffer:target_device_buffer->handle()
- destinationOffset:target_offset
- size:length];
-
- return OkStatus();
-}
-
-Status MetalCommandBuffer::PushConstants(iree_hal_executable_layout_t* executable_layout, size_t offset,
- absl::Span<const uint32_t> values) {
- IREE_TRACE_SCOPE0("MetalCommandBuffer::PushConstants");
- return UnimplementedErrorBuilder(IREE_LOC) << "MetalCommandBuffer::PushConstants";
-}
-
-Status MetalCommandBuffer::PushDescriptorSet(iree_hal_executable_layout_t* executable_layout, int32_t set,
- absl::Span<const iree_hal_descriptor_set_binding_t> bindings) {
- IREE_TRACE_SCOPE0("MetalCommandBuffer::PushDescriptorSet");
- if (set != 0) {
- return UnimplementedErrorBuilder(IREE_LOC)
- << "MetalCommandBuffer::PushDescriptorSet with set number > 0";
- }
- auto& push_state = pipeline_state_objects_[executable_layout].push_states[set];
- push_state.resource_bindings.assign(bindings.begin(), bindings.end());
- return OkStatus();
-}
-
-Status MetalCommandBuffer::BindDescriptorSet(iree_hal_executable_layout_t* executable_layout, int32_t set,
- iree_hal_descriptor_set_t* descriptor_set,
- absl::Span<const iree_device_size_t> dynamic_offsets) {
- IREE_TRACE_SCOPE0("MetalCommandBuffer::BindDescriptorSet");
- if (set != 0) {
- return UnimplementedErrorBuilder(IREE_LOC)
- << "MetalCommandBuffer::BindDescriptorSet with set number > 0";
- }
- if (!dynamic_offsets.empty()) {
- return UnimplementedErrorBuilder(IREE_LOC)
- << "MetalCommandBuffer::BindDescriptorSet with dynamic offsets";
- }
- pipeline_state_objects_[executable_layout].bind_states[set].descriptor_set = descriptor_set;
- return OkStatus();
-}
-
-Status MetalCommandBuffer::Dispatch(iree_hal_executable_t* executable, int32_t entry_point,
- std::array<uint32_t, 3> workgroups) {
- IREE_TRACE_SCOPE0("MetalCommandBuffer::Dispatch");
- IREE_DVLOG(2) << "MetalCommandBuffer::Dispatch";
-
- auto* kernel_library = static_cast<MetalKernelLibrary*>(executable);
- IREE_ASSIGN_OR_RETURN(auto metal_kernel, kernel_library->GetKernelForEntryPoint(entry_point));
- IREE_ASSIGN_OR_RETURN(auto metal_pso, kernel_library->GetPipelineStateForEntryPoint(entry_point));
- IREE_ASSIGN_OR_RETURN(auto workgroup_size,
- kernel_library->GetThreadgroupSizeForEntryPoint(entry_point));
-
- id<MTLComputeCommandEncoder> compute_encoder = GetOrBeginComputeEncoder();
- [compute_encoder setComputePipelineState:metal_pso];
-
- // TODO(antiagainst): only update the PSO for the current executable.
- for (const auto& pso_kv : pipeline_state_objects_) {
- const auto* pipeline_layout = static_cast<MetalPipelineArgumentBufferLayout*>(pso_kv.first);
-
- const auto& pso = pso_kv.second;
- if (pso.push_states.size() > 1) {
- return UnimplementedErrorBuilder(IREE_LOC)
- << "MetalCommandBuffer::Dispatch with more than one push descriptor sets";
- }
- if (!pso.bind_states.empty()) {
- return UnimplementedErrorBuilder(IREE_LOC)
- << "MetalCommandBuffer::Dispatch with bound descriptor sets";
- }
- if (!pso.constant_states.empty()) {
- return UnimplementedErrorBuilder(IREE_LOC)
- << "MetalCommandBuffer::Dispatch with push constants";
- }
-
- IREE_DVLOG(3) << "Encoding push descriptors..";
- for (const auto& push_kv : pso.push_states) {
- uint32_t set_number = push_kv.first;
- const PipelineStateObject::PushState& push_state = push_kv.second;
- IREE_DVLOG(3) << " For set #" << set_number;
-
- id<MTLArgumentEncoder> argument_encoder =
- [metal_kernel newArgumentEncoderWithBufferIndex:set_number]; // retained
- argument_encoder.label = @"IREE MetalCommandBuffer::Dispatch ArgumentEncoder";
- if (!argument_encoder) {
- return InvalidArgumentErrorBuilder(IREE_LOC)
- << "Buffer index #" << set_number << " is not an argument buffer";
- }
-
- __block id<MTLBuffer> argument_buffer =
- [metal_handle_.device newBufferWithLength:argument_encoder.encodedLength
- options:MTLResourceStorageModeShared]; // retained
- argument_encoder.label = @"IREE MetalCommandBuffer::Dispatch ArgumentBuffer";
- if (!argument_buffer) {
- return InternalErrorBuilder(IREE_LOC)
- << "Failed to create argument buffer with length=" << argument_encoder.encodedLength;
- }
- [metal_handle_ addCompletedHandler:^(id<MTLCommandBuffer>) {
- [argument_buffer release];
- [argument_encoder release];
- }];
-
- [argument_encoder setArgumentBuffer:argument_buffer offset:0];
-
- for (const auto& resource_binding : push_state.resource_bindings) {
-
- if (resource_binding.length != IREE_WHOLE_BUFFER &&
- resource_binding.length != resource_binding.buffer->allocation_size()) {
- return UnimplementedErrorBuilder(IREE_LOC)
- << "MetalCommandBuffer::Dispatch with sub-buffer";
- }
-
- [argument_encoder setBuffer:iree_hal_metal_buffer_handle(iree_hal_buffer_allocated_buffer(resource_binding.buffer))
- offset:resource_binding.offset
- atIndex:resource_binding.binding];
-
- const auto* set_layout = pipeline_layout->set_layouts()[set_number];
- const auto* layout_binding = set_layout->GetBindingForIndex(resource_binding.binding);
- if (!layout_binding) {
- return InvalidArgumentErrorBuilder(IREE_LOC)
- << "Cannot find binding #" << resource_binding.binding
- << " in argument buffer layout";
- }
- [compute_encoder useResource:buffer->handle()
- usage:ConvertResourceUsage(layout_binding->access)];
- }
-
- [compute_encoder setBuffer:argument_buffer offset:0 atIndex:set_number];
- }
- }
-
- IREE_DVLOG(2) << "Dispatch workgroup count: (" << workgroups[0] << ", " << workgroups[1] << ", "
- << workgroups[2] << "), workgroup size: (" << workgroup_size.x << ", "
- << workgroup_size.y << ", " << workgroup_size.z << ")";
- [compute_encoder
- dispatchThreadgroups:MTLSizeMake(workgroups[0], workgroups[1], workgroups[2])
- threadsPerThreadgroup:MTLSizeMake(workgroup_size.x, workgroup_size.y, workgroup_size.z)];
-
- return OkStatus();
-}
-
-Status MetalCommandBuffer::DispatchIndirect(iree_hal_executable_t* executable, int32_t entry_point,
- iree_hal_buffer_t* workgroups_buffer,
- iree_device_size_t workgroups_offset) {
- IREE_TRACE_SCOPE0("MetalCommandBuffer::DispatchIndirect");
- return UnimplementedErrorBuilder(IREE_LOC) << "MetalCommandBuffer::DispatchIndirect";
-}
-
-} // namespace metal
-} // namespace hal
-} // namespace iree
diff --git a/iree/hal/metal/metal_command_queue.h b/iree/hal/metal/metal_command_queue.h
deleted file mode 100644
index ef4d64c..0000000
--- a/iree/hal/metal/metal_command_queue.h
+++ /dev/null
@@ -1,54 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#ifndef IREE_HAL_METAL_METAL_COMMAND_QUEUE_H_
-#define IREE_HAL_METAL_METAL_COMMAND_QUEUE_H_
-
-#import <Metal/Metal.h>
-
-#include "iree/base/arena.h"
-#include "iree/base/status.h"
-#include "iree/base/time.h"
-#include "iree/hal/api.h"
-
-namespace iree {
-namespace hal {
-namespace metal {
-
-// A command queue implementation for Metal that directly wraps a
-// MTLCommandQueue.
-//
-// Thread-safe.
-class MetalCommandQueue final : public CommandQueue {
- public:
- MetalCommandQueue(std::string name,
- iree_hal_command_category_t supported_categories,
- id<MTLCommandQueue> queue);
- ~MetalCommandQueue() override;
-
- id<MTLCommandQueue> handle() const { return metal_handle_; }
-
- Status Submit(absl::Span<const SubmissionBatch> batches) override;
-
- Status WaitIdle(Time deadline_ns) override;
-
- private:
- id<MTLCommandQueue> metal_handle_;
-};
-
-} // namespace metal
-} // namespace hal
-} // namespace iree
-
-#endif // IREE_HAL_METAL_METAL_COMMAND_QUEUE_H_
diff --git a/iree/hal/metal/metal_command_queue.mm b/iree/hal/metal/metal_command_queue.mm
deleted file mode 100644
index eaf29fb..0000000
--- a/iree/hal/metal/metal_command_queue.mm
+++ /dev/null
@@ -1,100 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#include "iree/hal/metal/metal_command_queue.h"
-
-#include "iree/base/status.h"
-#include "iree/base/tracing.h"
-#include "iree/hal/metal/dispatch_time_util.h"
-#include "iree/hal/metal/metal_command_buffer.h"
-#include "iree/hal/metal/metal_shared_event.h"
-
-namespace iree {
-namespace hal {
-namespace metal {
-
-MetalCommandQueue::MetalCommandQueue(std::string name, iree_hal_command_category_t supported_categories,
- id<MTLCommandQueue> queue)
- : CommandQueue(std::move(name), supported_categories), metal_handle_([queue retain]) {
- metal_handle_.label = @"IREE MetalQueue";
-}
-
-MetalCommandQueue::~MetalCommandQueue() { [metal_handle_ release]; }
-
-Status MetalCommandQueue::Submit(absl::Span<const SubmissionBatch> batches) {
- IREE_TRACE_SCOPE0("MetalCommandQueue::Submit");
- for (const auto& batch : batches) {
- @autoreleasepool {
- // Wait for semaphores blocking this batch.
- if (!batch.wait_semaphores.empty()) {
- id<MTLCommandBuffer> wait_buffer = [metal_handle_ commandBufferWithUnretainedReferences];
- wait_buffer.label = @"IREE MetalCommandQueue::Submit Wait Semaphore CommandBuffer";
-
- for (const auto& semaphore : batch.wait_semaphores) {
- auto* event = static_cast<MetalSharedEvent*>(semaphore.semaphore);
- [wait_buffer encodeWaitForEvent:event->handle() value:semaphore.value];
- }
- [wait_buffer commit];
- }
-
- // Commit command buffers to the queue.
- for (const auto* command_buffer : batch.command_buffers) {
- const auto* cmdbuf = static_cast<const MetalCommandBuffer*>(command_buffer);
- [cmdbuf->handle() commit];
- }
-
- // Signal semaphores advanced by this batch.
- if (!batch.signal_semaphores.empty()) {
- id<MTLCommandBuffer> signal_buffer = [metal_handle_ commandBufferWithUnretainedReferences];
- signal_buffer.label = @"IREE MetalCommandQueue::Submit Signal Semaphore CommandBuffer";
-
- for (const auto& semaphore : batch.signal_semaphores) {
- auto* event = static_cast<MetalSharedEvent*>(semaphore.semaphore);
- [signal_buffer encodeSignalEvent:event->handle() value:semaphore.value];
- }
- [signal_buffer commit];
- }
- }
- }
- return OkStatus();
-}
-
-Status MetalCommandQueue::WaitIdle(Time deadline_ns) {
- IREE_TRACE_SCOPE0("MetalCommandQueue::WaitIdle");
-
- dispatch_time_t timeout = DeadlineToDispatchTime(deadline_ns);
-
- // Submit an empty command buffer and wait for it to complete. That will indicate all previous
- // work has completed too.
- @autoreleasepool {
- id<MTLCommandBuffer> comand_buffer = [metal_handle_ commandBufferWithUnretainedReferences];
- comand_buffer.label = @"IREE MetalCommandQueue::WaitIdle Command Buffer";
- __block dispatch_semaphore_t work_done = dispatch_semaphore_create(0);
- [comand_buffer addCompletedHandler:^(id<MTLCommandBuffer>) {
- dispatch_semaphore_signal(work_done);
- }];
- [comand_buffer commit];
- long timed_out = dispatch_semaphore_wait(work_done, timeout);
- dispatch_release(work_done);
- if (timed_out) {
- return DeadlineExceededErrorBuilder(IREE_LOC)
- << "Deadline exceeded waiting for dispatch_semaphore_t";
- }
- return OkStatus();
- }
-}
-
-} // namespace metal
-} // namespace hal
-} // namespace iree
diff --git a/iree/hal/metal/metal_device.h b/iree/hal/metal/metal_device.h
deleted file mode 100644
index aea32ce..0000000
--- a/iree/hal/metal/metal_device.h
+++ /dev/null
@@ -1,108 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#ifndef IREE_HAL_METAL_METAL_DEVICE_H_
-#define IREE_HAL_METAL_METAL_DEVICE_H_
-
-#import <Metal/Metal.h>
-
-#include <memory>
-
-#include "absl/types/span.h"
-#include "iree/base/memory.h"
-#include "iree/hal/cc/device.h"
-#include "iree/hal/cc/driver.h"
-
-namespace iree {
-namespace hal {
-namespace metal {
-
-// A device implementation for Metal that directly wraps a MTLDevice.
-class MetalDevice final : public Device {
- public:
- // Creates a device that retains the underlying Metal GPU device.
- // The iree_hal_device_id_t in |device_info| is expected to be an
- // id<MTLDevice>.
- static StatusOr<ref_ptr<MetalDevice>> Create(ref_ptr<Driver> driver,
- const DeviceInfo& device_info);
-
- ~MetalDevice() override;
-
- Allocator* allocator() const override { return allocator_.get(); }
-
- Status CreateExecutableCache(
- iree_string_view_t identifier,
- iree_hal_executable_cache_t** out_executable_cache) override;
-
- Status CreateDescriptorSetLayout(
- iree_hal_descriptor_set_layout_usage_type_t usage_type,
- absl::Span<const iree_hal_descriptor_set_layout_binding_t> bindings,
- iree_hal_descriptor_set_layout_t** out_descriptor_set_layout) override;
-
- Status CreateExecutableLayout(
- absl::Span<iree_hal_descriptor_set_layout_t*> set_layouts,
- size_t push_constants,
- iree_hal_executable_layout_t** out_executable_layout) override;
-
- Status CreateDescriptorSet(
- iree_hal_descriptor_set_layout_t* set_layout,
- absl::Span<const iree_hal_descriptor_set_binding_t> bindings,
- iree_hal_descriptor_set_t** out_descriptor_set) override;
-
- Status CreateCommandBuffer(
- iree_hal_command_buffer_mode_t mode,
- iree_hal_command_category_t command_categories,
- iree_hal_command_buffer_t** out_command_buffer) override;
-
- Status CreateEvent(iree_hal_event_t** out_event) override;
-
- Status CreateSemaphore(uint64_t initial_value,
- iree_hal_semaphore_t** out_semaphore) override;
- Status WaitAllSemaphores(const iree_hal_semaphore_list_t* semaphore_list,
- iree_time_t deadline_ns) override;
- StatusOr<int> WaitAnySemaphore(
- const iree_hal_semaphore_list_t* semaphore_list,
- iree_time_t deadline_ns) override;
-
- Status WaitIdle(iree_time_t deadline_ns) override;
-
- private:
- MetalDevice(ref_ptr<Driver> driver, const DeviceInfo& device_info);
-
- ref_ptr<Driver> driver_;
- id<MTLDevice> metal_handle_;
-
- std::unique_ptr<Allocator> allocator_;
-
- // Metal does not have clear graphics/dispatch/transfer queue distinction like
- // Vulkan; one just use the same newCommandQueue() API call on MTLDevice to
- // get command queues. Command encoders differ for different categories of
- // commands though. We expose one queue here for everything. This can be
- // changed later if more queues prove to be useful.
-
- std::unique_ptr<CommandQueue> command_queue_;
- mutable CommandQueue* common_queue_ = nullptr;
-
- // A dispatch queue and associated event listener for running Objective-C
- // blocks. This is typically used to wake up threads waiting on some HAL
- // semaphore.
- dispatch_queue_t wait_notifier_;
- MTLSharedEventListener* event_listener_;
-};
-
-} // namespace metal
-} // namespace hal
-} // namespace iree
-
-#endif // IREE_HAL_METAL_METAL_DEVICE_H_
diff --git a/iree/hal/metal/metal_device.mm b/iree/hal/metal/metal_device.mm
deleted file mode 100644
index 58e0000..0000000
--- a/iree/hal/metal/metal_device.mm
+++ /dev/null
@@ -1,210 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#include "iree/hal/metal/metal_device.h"
-
-#include "absl/memory/memory.h"
-#include "absl/strings/str_cat.h"
-#include "iree/base/status.h"
-#include "iree/base/time.h"
-#include "iree/base/tracing.h"
-#include "iree/hal/metal/dispatch_time_util.h"
-#include "iree/hal/metal/metal_capture_manager.h"
-#include "iree/hal/metal/metal_command_buffer.h"
-#include "iree/hal/metal/metal_command_queue.h"
-#include "iree/hal/metal/metal_direct_allocator.h"
-#include "iree/hal/metal/metal_pipeline_argument_buffer.h"
-#include "iree/hal/metal/metal_pipeline_cache.h"
-#include "iree/hal/metal/metal_shared_event.h"
-
-namespace iree {
-namespace hal {
-namespace metal {
-
-// static
-StatusOr<ref_ptr<MetalDevice>> MetalDevice::Create(ref_ptr<Driver> driver,
- const DeviceInfo& device_info,
- DebugCaptureManager* debug_capture_manager) {
- return assign_ref(new MetalDevice(std::move(driver), device_info, debug_capture_manager));
-}
-
-MetalDevice::MetalDevice(ref_ptr<Driver> driver, const DeviceInfo& device_info,
- DebugCaptureManager* debug_capture_manager)
- : Device(device_info),
- driver_(std::move(driver)),
- metal_handle_([(__bridge id<MTLDevice>)device_info.device_id() retain]),
- debug_capture_manager_(debug_capture_manager) {
- IREE_TRACE_SCOPE0("MetalDevice::ctor");
-
- // Grab one queue for dispatch and transfer.
- std::string name = absl::StrCat(device_info.name(), ":queue");
- id<MTLCommandQueue> metal_queue = [metal_handle_ newCommandQueue]; // retained
-
- allocator_ = MetalDirectAllocator::Create(metal_handle_, metal_queue);
-
- if (debug_capture_manager_ && debug_capture_manager_->is_connected()) {
- // Record a capture covering the duration of this device lifetime.
- static_cast<MetalCaptureManager*>(debug_capture_manager_)->SetCaptureObject(metal_handle_);
- debug_capture_manager_->StartCapture();
- }
-
- command_queue_ = absl::make_unique<MetalCommandQueue>(
- name, IREE_HAL_COMMAND_CATEGORY_ANY, metal_queue);
- common_queue_ = command_queue_.get();
- // MetalCommandQueue retains by itself. Release here to avoid leaking.
- [metal_queue release];
-
- wait_notifier_ = dispatch_queue_create("com.google.iree.semaphore_wait_notifier", NULL);
- event_listener_ = [[MTLSharedEventListener alloc] initWithDispatchQueue:wait_notifier_];
-}
-
-MetalDevice::~MetalDevice() {
- IREE_TRACE_SCOPE0("MetalDevice::dtor");
-
- if (debug_capture_manager_ && debug_capture_manager_->is_capturing()) {
- debug_capture_manager_->StopCapture();
- }
-
- [event_listener_ release];
- dispatch_release(wait_notifier_);
-
- [metal_handle_ release];
-}
-
-Status MetalDevice::CreateExecutableCache(iree_string_view_t identifier,
- iree_hal_executable_cache_t** out_executable_cache) {
- IREE_TRACE_SCOPE0("MetalDevice::CreateExecutableCache");
- return make_ref<MetalPipelineCache>(metal_handle_);
-}
-
-Status MetalDevice::CreateDescriptorSetLayout(
- iree_hal_descriptor_set_layout_usage_type_t usage_type,
- absl::Span<const iree_hal_descriptor_set_layout_binding_t> bindings,
- iree_hal_descriptor_set_layout_t** out_descriptor_set_layout) {
- IREE_TRACE_SCOPE0("MetalDevice::CreateDescriptorSetLayout");
- return make_ref<MetalArgumentBufferLayout>(usage_type, bindings);
-}
-
-Status MetalDevice::CreateExecutableLayout(
- absl::Span<iree_hal_descriptor_set_layout_t*> set_layouts, size_t push_constants, iree_hal_executable_layout_t** out_executable_layout) {
- IREE_TRACE_SCOPE0("MetalDevice::CreateExecutableLayout");
- return make_ref<MetalPipelineArgumentBufferLayout>(set_layouts, push_constants);
-}
-
-Status MetalDevice::CreateDescriptorSet(
- iree_hal_descriptor_set_layout_t* set_layout, absl::Span<const iree_hal_descriptor_set_binding_t> bindings,
- iree_hal_descriptor_set_t** out_descriptor_set) {
- IREE_TRACE_SCOPE0("MetalDevice::CreateDescriptorSet");
- return make_ref<MetalArgumentBuffer>(static_cast<MetalArgumentBufferLayout*>(set_layout),
- bindings);
-}
-
-Status MetalDevice::CreateCommandBuffer(
- iree_hal_command_buffer_mode_t mode, iree_hal_command_category_t command_categories,
- iree_hal_command_buffer_t** out_command_buffer) {
- IREE_TRACE_SCOPE0("MetalDevice::CreateCommandBuffer");
- @autoreleasepool {
- StatusOr<ref_ptr<CommandBuffer>> command_buffer;
- // We use commandBufferWithUnretainedReferences here to be performant. This is okay becasue
- // IREE tracks the lifetime of various objects with the help from compilers.
- id<MTLCommandBuffer> cmdbuf = [static_cast<MetalCommandQueue*>(common_queue_)->handle()
- commandBufferWithUnretainedReferences];
- command_buffer = MetalCommandBuffer::Create(mode, command_categories, cmdbuf);
- return command_buffer;
- }
-}
-
-Status MetalDevice::CreateEvent(iree_hal_event_t** out_event) {
- IREE_TRACE_SCOPE0("MetalDevice::CreateEvent");
- return UnimplementedErrorBuilder(IREE_LOC) << "MetalDevice::CreateEvent";
-}
-
-Status MetalDevice::CreateSemaphore(uint64_t initial_value, iree_hal_semaphore_t** out_semaphore) {
- IREE_TRACE_SCOPE0("MetalDevice::CreateSemaphore");
- return MetalSharedEvent::Create(metal_handle_, event_listener_, initial_value);
-}
-
-Status MetalDevice::WaitAllSemaphores(const iree_hal_semaphore_list_t* semaphore_list,
- iree_time_t deadline_ns) {
- IREE_TRACE_SCOPE0("MetalDevice::WaitAllSemaphores");
- // Go through all MetalSharedEvents and wait on each of them given we need all of them to be
- // signaled anyway.
- for (int i = 0; i < semaphores.size(); ++i) {
- auto* semaphore = static_cast<MetalSharedEvent*>(semaphores[i].semaphore);
- IREE_RETURN_IF_ERROR(semaphore->Wait(semaphores[i].value, deadline_ns));
- }
- return OkStatus();
-}
-
-StatusOr<int> MetalDevice::WaitAnySemaphore(const iree_hal_semaphore_list_t* semaphore_list,
- iree_time_t deadline_ns) {
- IREE_TRACE_SCOPE0("MetalDevice::WaitAnySemaphore");
-
- if (semaphores.empty()) {
- return InvalidArgumentErrorBuilder(IREE_LOC) << "expected to have at least one semaphore";
- }
-
- // If there is just one semaphore, just wait on it.
- if (semaphores.size() == 1) {
- auto* semaphore = static_cast<MetalSharedEvent*>(semaphores[0].semaphore);
- IREE_RETURN_IF_ERROR(semaphore->Wait(semaphores[0].value, deadline_ns));
- return 0;
- }
-
- // Otherwise, we need to go down a more complicated path by registering listeners to all
- // MTLSharedEvents to notify us when at least one of them has done the work on GPU by signaling a
- // semaphore. The signaling will happen in a new dispatch queue; the current thread will wait on
- // the semaphore.
-
- dispatch_time_t timeout = DeadlineToDispatchTime(deadline_ns);
-
- // Store the handle as a __block variable to allow blocks accessing the same copy for the
- // semaphore handle on heap.
- // Use an initial value of zero so that any semaphore signal will unblock the wait.
- __block dispatch_semaphore_t work_done = dispatch_semaphore_create(0);
- // Also create a __block variable to store the index for the signaled semaphore.
- __block int signaled_index = 0;
-
- // The dispatch queue created in the above is a serial one. So even if multiple semaphores signal,
- // the semaphore signaling should be serialized.
- for (int i = 0; i < semaphores.size(); ++i) {
- auto* semaphore = static_cast<MetalSharedEvent*>(semaphores[i].semaphore);
- [semaphore->handle() notifyListener:event_listener_
- atValue:semaphores[i].value
- block:^(id<MTLSharedEvent>, uint64_t) {
- dispatch_semaphore_signal(work_done);
- // This should capture the *current* index for each semaphore.
- signaled_index = i;
- }];
- }
-
- long timed_out = dispatch_semaphore_wait(work_done, timeout);
-
- dispatch_release(work_done);
-
- if (timed_out) {
- return DeadlineExceededErrorBuilder(IREE_LOC)
- << "Deadline exceeded waiting for dispatch_semaphore_t";
- }
- return signaled_index;
-}
-
-Status MetalDevice::WaitIdle(iree_time_t deadline_ns) {
- IREE_TRACE_SCOPE0("MetalDevice::WaitIdle");
- return UnimplementedErrorBuilder(IREE_LOC) << "MetalDevice::WaitIdle";
-}
-
-} // namespace metal
-} // namespace hal
-} // namespace iree
diff --git a/iree/hal/metal/metal_direct_allocator.h b/iree/hal/metal/metal_direct_allocator.h
deleted file mode 100644
index 117d358..0000000
--- a/iree/hal/metal/metal_direct_allocator.h
+++ /dev/null
@@ -1,66 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#ifndef IREE_HAL_METAL_METAL_DIRECT_ALLOCATOR_H_
-#define IREE_HAL_METAL_METAL_DIRECT_ALLOCATOR_H_
-
-#import <Metal/Metal.h>
-
-#include <memory>
-
-#include "iree/base/status.h"
-#include "iree/hal/api.h"
-
-namespace iree {
-namespace hal {
-namespace metal {
-
-class MetalBuffer;
-
-// An allocator implementation for Metal that directly wraps a MTLDevice and
-// requests all allocations on the device. This is not of great performance,
-// but good for start.
-class MetalDirectAllocator final : public Allocator {
- public:
- static std::unique_ptr<MetalDirectAllocator> Create(
- id<MTLDevice> device, id<MTLCommandQueue> transfer_queue);
-
- ~MetalDirectAllocator() override;
-
- bool CanUseBufferLike(Allocator* source_allocator,
- iree_hal_memory_type_t memory_type,
- iree_hal_buffer_usage_t buffer_usage,
- iree_hal_buffer_usage_t intended_usage) const override;
-
- StatusOr<ref_ptr<Buffer>> Allocate(iree_hal_memory_type_t memory_type,
- iree_hal_buffer_usage_t buffer_usage,
- size_t allocation_size) override;
-
- private:
- explicit MetalDirectAllocator(id<MTLDevice> device,
- id<MTLCommandQueue> transfer_queue);
-
- StatusOr<ref_ptr<MetalBuffer>> AllocateInternal(
- iree_hal_memory_type_t memory_type, iree_hal_buffer_usage_t buffer_usage,
- iree_hal_memory_access_t allowed_access, size_t allocation_size);
-
- id<MTLDevice> metal_device_;
- id<MTLCommandQueue> metal_transfer_queue_;
-};
-
-} // namespace metal
-} // namespace hal
-} // namespace iree
-
-#endif // IREE_HAL_METAL_METAL_DIRECT_ALLOCATOR_H_
diff --git a/iree/hal/metal/metal_direct_allocator.mm b/iree/hal/metal/metal_direct_allocator.mm
deleted file mode 100644
index 833006d..0000000
--- a/iree/hal/metal/metal_direct_allocator.mm
+++ /dev/null
@@ -1,124 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#include "iree/hal/metal/metal_direct_allocator.h"
-
-#include "absl/memory/memory.h"
-#include "iree/base/status.h"
-#include "iree/base/tracing.h"
-#include "iree/hal/metal/metal_buffer.h"
-
-namespace iree {
-namespace hal {
-namespace metal {
-
-namespace {
-
-// Returns the proper Metal resource storage mode given the specific MemoryType.
-MTLResourceOptions SelectMTLResourceStorageMode(MemoryType memory_type) {
- // There are four MTLStorageMode:
- // * Managed: The CPU and GPU may maintain separate copies of the resource, and any changes
- // must be explicitly synchronized.
- // * Shared: The resource is stored in system memory and is accessible to both the CPU and
- // the GPU.
- // * Private: The resource can be accessed only by the GPU.
- // * Memoryless: The resource’s contents can be accessed only by the GPU and only exist
- // temporarily during a render pass.
- // macOS has all of the above; MTLStorageModeManaged is not available on iOS.
- //
- // The IREE HAL is modeled after Vulkan so it's quite explicit. For buffers visible to both
- // the host and the device, we would like to opt in with the explicit version
- // (MTLStorageManaged) when possible because it should be more performant: "In macOS,
- // there’s no difference in GPU performance between managed and private buffers." But for
- // iOS, MTLStorageShared should be good given we have a unified memory model there.
-
- if (iree_all_bits_set(memory_type, IREE_HAL_MEMORY_TYPE_DEVICE_LOCAL)) {
- if (iree_all_bits_set(memory_type, IREE_HAL_MEMORY_TYPE_HOST_VISIBLE)) {
- // Device-local, host-visible.
- // TODO(antiagainst): Enable using MTLResourceStorageModeManaged on macOS once we have
- // defined invalidate/flush C APIs and wired up their usage through the stack. At the
- // moment if we use MTLResourceStorageModeManaged, due to no proper invlidate/flush
- // actions, the kernel invocations' data read/write will not be properly synchronized.
- return MTLResourceStorageModeShared;
- } else {
- // Device-local only.
- return MTLResourceStorageModePrivate;
- }
- } else {
- if (iree_all_bits_set(memory_type, IREE_HAL_MEMORY_TYPE_DEVICE_VISIBLE)) {
- // Host-local, device-visible.
- return MTLResourceStorageModeShared;
- } else {
- // Host-local only.
- // TODO(antiagainst): we probably want to just use HostBuffer here.
- return MTLResourceStorageModeShared;
- }
- }
-}
-
-} // namespace
-
-// static
-std::unique_ptr<MetalDirectAllocator> MetalDirectAllocator::Create(
- id<MTLDevice> device, id<MTLCommandQueue> transfer_queue) {
- IREE_TRACE_SCOPE0("MetalDirectAllocator::Create");
- return absl::WrapUnique(new MetalDirectAllocator(device, transfer_queue));
-}
-
-MetalDirectAllocator::MetalDirectAllocator(id<MTLDevice> device, id<MTLCommandQueue> transfer_queue)
- : metal_device_([device retain]), metal_transfer_queue_([transfer_queue retain]) {}
-
-MetalDirectAllocator::~MetalDirectAllocator() {
- IREE_TRACE_SCOPE0("MetalDirectAllocator::dtor");
- [metal_transfer_queue_ release];
- [metal_device_ release];
-}
-
-bool MetalDirectAllocator::CanUseBufferLike(Allocator* source_allocator,
- iree_hal_memory_type_t memory_type,
- iree_hal_buffer_usage_t buffer_usage,
- iree_hal_buffer_usage_t intended_usage) const {
- // TODO(benvanik): ensure there is a memory type that can satisfy the request.
- return source_allocator == this;
-}
-
-StatusOr<ref_ptr<MetalBuffer>> MetalDirectAllocator::AllocateInternal(
- iree_hal_memory_type_t memory_type, iree_hal_buffer_usage_t buffer_usage,
- iree_hal_memory_access_t allowed_access, size_t allocation_size) {
- IREE_TRACE_SCOPE0("MetalDirectAllocator::AllocateInternal");
-
- MTLResourceOptions resource_options = SelectMTLResourceStorageMode(memory_type);
-
- // IREE is more explicit than Metal: it tracks various state by itself. There is no need
- // to incur Metal runtime overhead for hazard tracking.
- resource_options |= MTLResourceHazardTrackingModeUntracked;
-
- id<MTLBuffer> metal_buffer = [metal_device_ newBufferWithLength:allocation_size
- options:resource_options]; // retained
-
- return MetalBuffer::Create(
- this, memory_type, allowed_access, buffer_usage, allocation_size, /*byte_offset=*/0,
- /*byte_length=*/allocation_size, metal_buffer, metal_transfer_queue_);
-}
-
-StatusOr<ref_ptr<Buffer>> MetalDirectAllocator::Allocate(iree_hal_memory_type_t memory_type,
- iree_hal_buffer_usage_t buffer_usage,
- size_t allocation_size) {
- IREE_TRACE_SCOPE0("MetalDirectAllocator::Allocate");
- return AllocateInternal(memory_type, buffer_usage, IREE_HAL_MEMORY_ACCESS_ALL, allocation_size);
-}
-
-} // namespace metal
-} // namespace hal
-} // namespace iree
diff --git a/iree/hal/metal/metal_driver.h b/iree/hal/metal/metal_driver.h
deleted file mode 100644
index bb2603e..0000000
--- a/iree/hal/metal/metal_driver.h
+++ /dev/null
@@ -1,62 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#ifndef IREE_HAL_METAL_METAL_DRIVER_H_
-#define IREE_HAL_METAL_METAL_DRIVER_H_
-
-#include <memory>
-#include <string>
-
-#include "iree/hal/cc/driver.h"
-
-namespace iree {
-namespace hal {
-namespace metal {
-
-struct MetalDriverOptions {
- // Whether to enable Metal command capture.
- bool enable_capture;
- // The file to contain the Metal capture. Empty means capturing to Xcode.
- std::string capture_file;
-};
-
-// A pseudo Metal GPU driver which retains all available Metal GPU devices
-// during its lifetime.
-//
-// It uses the iree_hal_device_id_t to store the underlying id<MTLDevice>.
-class MetalDriver final : public Driver {
- public:
- static StatusOr<ref_ptr<MetalDriver>> Create(
- const MetalDriverOptions& options);
-
- ~MetalDriver() override;
-
- StatusOr<std::vector<DeviceInfo>> EnumerateAvailableDevices() override;
-
- StatusOr<ref_ptr<Device>> CreateDefaultDevice() override;
-
- StatusOr<ref_ptr<Device>> CreateDevice(
- iree_hal_device_id_t device_id) override;
-
- private:
- MetalDriver(std::vector<DeviceInfo> devices);
-
- std::vector<DeviceInfo> devices_;
-};
-
-} // namespace metal
-} // namespace hal
-} // namespace iree
-
-#endif // IREE_HAL_METAL_METAL_DRIVER_H_
diff --git a/iree/hal/metal/metal_driver.mm b/iree/hal/metal/metal_driver.mm
deleted file mode 100644
index b561a31..0000000
--- a/iree/hal/metal/metal_driver.mm
+++ /dev/null
@@ -1,118 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#include "iree/hal/metal/metal_driver.h"
-
-#import <Metal/Metal.h>
-
-#include "iree/base/status.h"
-#include "iree/base/tracing.h"
-#include "iree/hal/metal/metal_capture_manager.h"
-#include "iree/hal/metal/metal_device.h"
-
-namespace iree {
-namespace hal {
-namespace metal {
-
-namespace {
-
-// Returns an autoreleased array of available Metal GPU devices.
-NSArray<id<MTLDevice>>* GetAvailableMetalDevices() {
-#if defined(IREE_PLATFORM_MACOS)
- // For macOS, we might have more than one GPU devices.
- return [MTLCopyAllDevices() autorelease];
-#else
- // For other Apple platforms, we only have one GPU device.
- id<MTLDevice> device = [MTLCreateSystemDefaultDevice() autorelease];
- return [NSArray arrayWithObject:device];
-#endif
-}
-
-} // namespace
-
-// static
-StatusOr<ref_ptr<MetalDriver>> MetalDriver::Create(const MetalDriverOptions& options) {
- IREE_TRACE_SCOPE0("MetalDriver::Create");
-
- @autoreleasepool {
- NSArray<id<MTLDevice>>* devices = GetAvailableMetalDevices();
- if (devices == nil) {
- return UnavailableErrorBuilder(IREE_LOC) << "no Metal GPU devices available";
- }
-
- std::unique_ptr<MetalCaptureManager> metal_capture_manager;
- if (options.enable_capture) {
- IREE_ASSIGN_OR_RETURN(metal_capture_manager,
- MetalCaptureManager::Create(options.capture_file));
- IREE_RETURN_IF_ERROR(metal_capture_manager->Connect());
- }
-
- std::vector<DeviceInfo> device_infos;
- for (id<MTLDevice> device in devices) {
- std::string name = std::string([device.name UTF8String]);
- iree_hal_device_feature_t supported_features = IREE_HAL_DEVICE_FEATURE_NONE;
- iree_hal_device_id_t device_id = reinterpret_cast<iree_hal_device_id_t>((__bridge void*)device);
- device_infos.emplace_back("metal", std::move(name), supported_features, device_id);
- }
- return assign_ref(new MetalDriver(std::move(device_infos), std::move(metal_capture_manager)));
- }
-}
-
-MetalDriver::MetalDriver(std::vector<DeviceInfo> devices)
- : Driver("metal"),
- devices_(std::move(devices)) {
- // Retain all the retained Metal GPU devices.
- for (const auto& device : devices_) {
- [(__bridge id<MTLDevice>)device.device_id() retain];
- }
-}
-
-MetalDriver::~MetalDriver() {
- IREE_TRACE_SCOPE0("MetalDriver::dtor");
-
- // Release all the retained Metal GPU devices.
- for (const auto& device : devices_) {
- [(__bridge id<MTLDevice>)device.device_id() release];
- }
-}
-
-StatusOr<std::vector<DeviceInfo>> MetalDriver::EnumerateAvailableDevices() {
- IREE_TRACE_SCOPE0("MetalDriver::EnumerateAvailableDevices");
-
- return devices_;
-}
-
-StatusOr<ref_ptr<Device>> MetalDriver::CreateDefaultDevice() {
- IREE_TRACE_SCOPE0("MetalDriver::CreateDefaultDevice");
-
- if (devices_.empty()) {
- return UnavailableErrorBuilder(IREE_LOC) << "no Metal GPU devices available";
- }
- return CreateDevice(devices_.front().device_id());
-}
-
-StatusOr<ref_ptr<Device>> MetalDriver::CreateDevice(iree_hal_device_id_t device_id) {
- IREE_TRACE_SCOPE0("MetalDriver::CreateDevice");
-
- for (const DeviceInfo& info : devices_) {
- if (info.device_id() == device_id) {
- return MetalDevice::Create(add_ref(this), info);
- }
- }
- return InvalidArgumentErrorBuilder(IREE_LOC) << "unknown driver device id: " << device_id;
-}
-
-} // namespace metal
-} // namespace hal
-} // namespace iree
diff --git a/iree/hal/metal/metal_kernel_library.h b/iree/hal/metal/metal_kernel_library.h
deleted file mode 100644
index 6916ddc..0000000
--- a/iree/hal/metal/metal_kernel_library.h
+++ /dev/null
@@ -1,86 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#ifndef IREE_HAL_METAL_METAL_KERNEL_LIBRARY_H_
-#define IREE_HAL_METAL_METAL_KERNEL_LIBRARY_H_
-
-#import <Metal/Metal.h>
-
-#include <string>
-
-#include "absl/container/inlined_vector.h"
-#include "iree/base/status.h"
-#include "iree/hal/cc/executable_cache.h"
-
-// flatcc schemas:
-#include "iree/base/flatcc.h"
-#include "iree/schemas/metal_executable_def_builder.h"
-#include "iree/schemas/metal_executable_def_reader.h"
-#include "iree/schemas/metal_executable_def_verifier.h"
-
-namespace iree {
-namespace hal {
-namespace metal {
-
-// An executable implementation for Metal that wraps MTLLibrary and MTLFunction.
-//
-// Metal represents compute kernels as MTLFunctions. MTLLibrary is just an
-// allocation of MTLFunctions. One creates a MTLComputePipelineState from a
-// MTLFunction and uses the pipeline state for creating compute pipelines.
-// This class bundles all the necesary Metal objects for getting pipeline state
-// objects for a compute kernel.
-class MetalKernelLibrary final : public Executable {
- public:
- static StatusOr<ref_ptr<MetalKernelLibrary>> Create(
- id<MTLDevice> device, iree_hal_executable_caching_mode_t mode,
- iree_const_byte_span_t executable_data);
- ~MetalKernelLibrary() override;
-
- // Returns the MTLFunction for the entry point with the given |ordinal|.
- StatusOr<id<MTLFunction>> GetKernelForEntryPoint(int ordinal) const;
-
- // Returns the threadgroup size for the entry point with the given |ordinal|.
- StatusOr<iree_MetalThreadgroupSize_t> GetThreadgroupSizeForEntryPoint(
- int ordinal) const;
-
- // Returns the pipeline state object for the entry point with the given
- // |ordinal|.
- StatusOr<id<MTLComputePipelineState>> GetPipelineStateForEntryPoint(
- int ordinal) const;
-
- private:
- struct KernelObjects {
- id<MTLFunction> function;
- iree_MetalThreadgroupSize_t threadgroup_size;
- // Baked pipeline state object.
- id<MTLComputePipelineState> pipeline_state;
- };
-
- // Creates a MetalKernelLibrary assuming all Metal objects are already
- // retained before passing in.
- MetalKernelLibrary(id<MTLDevice> device,
- absl::InlinedVector<id<MTLLibrary>, 4> libraries,
- absl::InlinedVector<KernelObjects, 4> kernel_objects);
-
- id<MTLDevice> device_;
-
- absl::InlinedVector<id<MTLLibrary>, 4> libraries_;
- absl::InlinedVector<KernelObjects, 4> kernel_objects_;
-};
-
-} // namespace metal
-} // namespace hal
-} // namespace iree
-
-#endif // IREE_HAL_METAL_METAL_KERNEL_LIBRARY_H_
diff --git a/iree/hal/metal/metal_kernel_library.mm b/iree/hal/metal/metal_kernel_library.mm
deleted file mode 100644
index dc1f733..0000000
--- a/iree/hal/metal/metal_kernel_library.mm
+++ /dev/null
@@ -1,225 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#include "iree/hal/metal/metal_kernel_library.h"
-
-#include "iree/base/memory.h"
-#include "iree/base/status.h"
-#include "iree/base/tracing.h"
-
-// NOTE: starting to port this to ObjC.
-
-// Verifies the structure of the flatbuffer so that we can avoid doing so during
-// runtime. There are still some conditions we must be aware of (such as omitted
-// names on functions with internal linkage), however we shouldn't need to
-// bounds check anything within the flatbuffer after this succeeds.
-static iree_status_t iree_hal_metal_executable_flatbuffer_verify(
- iree_const_byte_span_t flatbuffer_data) {
- if (!flatbuffer_data.data || flatbuffer_data.data_length < 16) {
- return iree_make_status(IREE_STATUS_INVALID_ARGUMENT,
- "flatbuffer data is not present or less than 16 bytes (%zu total)",
- flatbuffer_data.data_length);
- }
-
- // Run flatcc generated verification. This ensures all pointers are in-bounds
- // and that we can safely walk the file, but not that the actual contents of
- // the flatbuffer meet our expectations.
- int verify_ret =
- iree_MetalExecutableDef_verify_as_root(flatbuffer_data.data, flatbuffer_data.data_length);
- if (verify_ret != flatcc_verify_ok) {
- return iree_make_status(IREE_STATUS_INVALID_ARGUMENT, "flatbuffer verification failed: %s",
- flatcc_verify_error_string(verify_ret));
- }
-
- iree_MetalExecutableDef_table_t executable_def =
- iree_MetalExecutableDef_as_root(flatbuffer_data.data);
-
- flatbuffers_string_vec_t entry_points_vec =
- iree_MetalExecutableDef_entry_points_get(executable_def);
- size_t entry_point_count = flatbuffers_string_vec_len(entry_points_vec);
- for (size_t i = 0; i < entry_point_count; ++i) {
- if (!flatbuffers_string_len(flatbuffers_string_vec_at(entry_points_vec, i))) {
- return iree_make_status(IREE_STATUS_INVALID_ARGUMENT,
- "executable entry point %zu has no name", i);
- }
- }
-
- iree_MetalThreadgroupSize_vec_t threadgroup_sizes_vec =
- iree_MetalExecutableDef_threadgroup_sizes(executable_def);
- size_t threadgroup_size_count = iree_MetalThreadgroupSize_vec_len(threadgroup_sizes_vec);
- if (!threadgroup_size_count) {
- return InvalidArgumentErrorBuilder(IREE_LOC) << "No threadgroup sizes present";
- }
-
- flatbuffers_string_vec_t shader_sources_vec =
- iree_MetalExecutableDef_shader_sources_get(executable_def);
- size_t shader_source_count = flatbuffers_string_vec_len(shader_sources_vec);
- for (size_t i = 0; i < shader_source_count; ++i) {
- if (!flatbuffers_string_len(flatbuffers_string_vec_at(shader_sources_vec, i))) {
- return iree_make_status(IREE_STATUS_INVALID_ARGUMENT, "executable shader source %zu is empty",
- i);
- }
- }
-
- if (entry_point_count != threadgroup_size_count || entry_point_count != shader_source_count) {
- return iree_make_status(IREE_STATUS_INVALID_ARGUMENT,
- "mismatch among the numbers of entry points (%zu), thread group sizes "
- "(%zu), and source strings (%zu)",
- entry_point_count, threadgroup_size_count, shader_source_count);
- }
-
- return iree_ok_status();
-}
-
-namespace iree {
-namespace hal {
-namespace metal {
-
-// static
-StatusOr<ref_ptr<MetalKernelLibrary>> MetalKernelLibrary::Create(id<MTLDevice> device,
- iree_hal_executable_caching_mode_t mode,
- iree_const_byte_span_t executable_data) {
- IREE_TRACE_SCOPE0("MetalKernelLibrary::Create");
-
- // Verify and fetch the executable flatbuffer wrapper.
- iree_const_byte_span_t executable_data =
- iree_make_const_byte_span(spec.executable_data.data(), spec.executable_data.size());
- IREE_RETURN_IF_ERROR(iree_hal_metal_executable_flatbuffer_verify(executable_data));
- iree_MetalExecutableDef_table_t executable_def =
- iree_MetalExecutableDef_as_root(executable_data.data);
-
- flatbuffers_string_vec_t entry_points_vec =
- iree_MetalExecutableDef_entry_points_get(executable_def);
- iree_MetalThreadgroupSize_vec_t threadgroup_sizes_vec =
- iree_MetalExecutableDef_threadgroup_sizes(executable_def);
- flatbuffers_string_vec_t shader_sources_vec =
- iree_MetalExecutableDef_shader_sources_get(executable_def);
-
- // Compile each MSL source string into a MTLLibrary and get the MTLFunction for the entry point to
- // build the pipeline state object.
-
- absl::InlinedVector<id<MTLLibrary>, 4> libraries;
- absl::InlinedVector<KernelObjects, 4> kernel_objects;
-
- MTLCompileOptions* msl_compile_options = [MTLCompileOptions new];
- msl_compile_options.languageVersion = MTLLanguageVersion2_0;
-
- auto cleanup = MakeCleanup([&]() {
- for (const auto& kernel : kernel_objects) {
- [kernel.pipeline_state release];
- [kernel.function release];
- }
- for (id<MTLLibrary> library : libraries) [library release];
- [msl_compile_options release];
- });
-
- // TODO(antiagainst): We are performing synchronous compilation at runtime here. This is good for
- // debugging purposes but bad for performance. Enable offline compilation and make that as the
- // default.
-
- for (size_t entry_ordinal = 0; entry_ordinal < flatbuffers_string_vec_len(shader_sources_vec);
- ++entry_ordinal) {
- flatbuffers_string_t entry_point = flatbuffers_string_vec_at(entry_points_vec, entry_ordinal);
- @autoreleasepool {
- NSError* error = nil;
-
- NSString* shader_source =
- [NSString stringWithCString:flatbuffers_string_vec_at(shader_sources_vec, entry_ordinal)
- encoding:[NSString defaultCStringEncoding]];
- id<MTLLibrary> library = [device newLibraryWithSource:shader_source
- options:msl_compile_options
- error:&error];
- if (!library) {
- NSLog(@"Failed to create MTLLibrary: %@", error);
-#ifndef NDEBUG
- NSLog(@"Original MSL source: %@", shader_source);
-#endif
- return InvalidArgumentErrorBuilder(IREE_LOC) << "Invalid MSL source";
- }
- libraries.push_back(library);
-
- id<MTLFunction> function = [library
- newFunctionWithName:[NSString stringWithCString:entry_point
- encoding:[NSString defaultCStringEncoding]]];
- if (!function) {
- NSLog(@"Failed to create MTLFunction");
-#ifndef NDEBUG
- NSLog(@"Original MSL source: %@", shader_source);
-#endif
- return InvalidArgumentErrorBuilder(IREE_LOC)
- << "Cannot find entry point '" << entry_point << "' in shader source index "
- << entry_ordinal;
- }
-
- id<MTLComputePipelineState> pso = [device newComputePipelineStateWithFunction:function
- error:&error];
- if (!pso) {
- NSLog(@"Failed to create MTLComputePipelineState: %@", error);
-#ifndef NDEBUG
- NSLog(@"Original MSL source: %@", shader_source);
-#endif
- return InvalidArgumentErrorBuilder(IREE_LOC) << "Invalid MSL source";
- }
-
- kernel_objects.push_back(
- KernelObjects{function, {static_cast<uint32_t>(iree_MetalThreadgroupSize__size())}, pso});
- }
- }
-
- return assign_ref(
- new MetalKernelLibrary([device retain], std::move(libraries), std::move(kernel_objects)));
-}
-
-MetalKernelLibrary::MetalKernelLibrary(id<MTLDevice> device,
- absl::InlinedVector<id<MTLLibrary>, 4> libraries,
- absl::InlinedVector<KernelObjects, 4> kernel_objects)
- : device_(device),
- libraries_(std::move(libraries)),
- kernel_objects_(std::move(kernel_objects)) {}
-
-MetalKernelLibrary::~MetalKernelLibrary() {
- IREE_TRACE_SCOPE0("MetalKernelLibrary::dtor");
- for (const auto& kernel : kernel_objects_) {
- [kernel.pipeline_state release];
- [kernel.function release];
- }
- for (id<MTLLibrary> library : libraries_) [library release];
-}
-
-StatusOr<id<MTLFunction>> MetalKernelLibrary::GetKernelForEntryPoint(int ordinal) const {
- if (ordinal < 0 || ordinal >= kernel_objects_.size()) {
- return OutOfRangeErrorBuilder(IREE_LOC) << "Invalid entry point ordinal: " << ordinal;
- }
- return kernel_objects_[ordinal].function;
-}
-
-StatusOr<iree_MetalThreadgroupSize_t> MetalKernelLibrary::GetThreadgroupSizeForEntryPoint(
- int ordinal) const {
- if (ordinal < 0 || ordinal >= kernel_objects_.size()) {
- return OutOfRangeErrorBuilder(IREE_LOC) << "Invalid entry point ordinal: " << ordinal;
- }
- return kernel_objects_[ordinal].threadgroup_size;
-}
-
-StatusOr<id<MTLComputePipelineState>> MetalKernelLibrary::GetPipelineStateForEntryPoint(
- int ordinal) const {
- if (ordinal < 0 || ordinal >= kernel_objects_.size()) {
- return OutOfRangeErrorBuilder(IREE_LOC) << "Invalid entry point ordinal: " << ordinal;
- }
- return kernel_objects_[ordinal].pipeline_state;
-}
-
-} // namespace metal
-} // namespace hal
-} // namespace iree
diff --git a/iree/hal/metal/metal_pipeline_argument_buffer.cc b/iree/hal/metal/metal_pipeline_argument_buffer.cc
deleted file mode 100644
index 1879a81..0000000
--- a/iree/hal/metal/metal_pipeline_argument_buffer.cc
+++ /dev/null
@@ -1,61 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#include "iree/hal/metal/metal_pipeline_argument_buffer.h"
-
-#include "absl/strings/str_cat.h"
-#include "absl/strings/str_join.h"
-
-namespace iree {
-namespace hal {
-namespace metal {
-
-MetalArgumentBufferLayout::MetalArgumentBufferLayout(
- iree_hal_descriptor_set_layout_usage_type_t usage_type,
- absl::Span<const iree_hal_descriptor_set_layout_binding_t> bindings)
- : usage_type_(usage_type), bindings_(bindings.begin(), bindings.end()) {}
-
-const iree_hal_descriptor_set_layout_binding_t*
-MetalArgumentBufferLayout::GetBindingForIndex(int index) const {
- for (const auto& binding : bindings_) {
- if (binding.binding == index) return &binding;
- }
- return nullptr;
-}
-
-MetalPipelineArgumentBufferLayout::MetalPipelineArgumentBufferLayout(
- absl::Span<DescriptorSetLayout* const> set_layouts, size_t push_constants)
- : set_layouts_(set_layouts.size()), push_constants_(push_constants) {
- for (int i = 0; i < set_layouts.size(); ++i) {
- set_layouts_[i] = static_cast<MetalArgumentBufferLayout*>(set_layouts[i]);
- set_layouts_[i]->AddReference();
- }
-}
-
-MetalPipelineArgumentBufferLayout::~MetalPipelineArgumentBufferLayout() {
- for (auto* layout : set_layouts_) layout->ReleaseReference();
-}
-
-MetalArgumentBuffer::MetalArgumentBuffer(
- MetalArgumentBufferLayout* layout,
- absl::Span<const iree_hal_descriptor_set_binding_t> resources)
- : layout_(layout), resources_(resources.begin(), resources.end()) {
- layout_->AddReference();
-}
-
-MetalArgumentBuffer::~MetalArgumentBuffer() { layout_->ReleaseReference(); }
-
-} // namespace metal
-} // namespace hal
-} // namespace iree
diff --git a/iree/hal/metal/metal_pipeline_argument_buffer.h b/iree/hal/metal/metal_pipeline_argument_buffer.h
deleted file mode 100644
index 0d5456a..0000000
--- a/iree/hal/metal/metal_pipeline_argument_buffer.h
+++ /dev/null
@@ -1,80 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#ifndef IREE_HAL_METAL_METAL_PIPELINE_ARGUMENT_BUFFER_H_
-#define IREE_HAL_METAL_METAL_PIPELINE_ARGUMENT_BUFFER_H_
-
-#include <string>
-
-#include "absl/container/inlined_vector.h"
-#include "absl/types/span.h"
-#include "iree/hal/cc/descriptor_set.h"
-#include "iree/hal/cc/descriptor_set_layout.h"
-#include "iree/hal/cc/executable_layout.h"
-
-// Metal implementaion classes for resource descriptor related interfaces.
-//
-// See docs/design_docs/metal_hal_driver.md#resource-descriptors for more
-// details.
-
-namespace iree {
-namespace hal {
-namespace metal {
-
-class MetalArgumentBufferLayout final : public DescriptorSetLayout {
- public:
- MetalArgumentBufferLayout(UsageType usage_type,
- absl::Span<const Binding> bindings);
- ~MetalArgumentBufferLayout() override = default;
-
- absl::Span<const Binding> bindings() const { return bindings_; }
- const Binding* GetBindingForIndex(int index) const;
-
- private:
- UsageType usage_type_;
- absl::InlinedVector<Binding, 8> bindings_;
-};
-
-class MetalPipelineArgumentBufferLayout final : public ExecutableLayout {
- public:
- MetalPipelineArgumentBufferLayout(
- absl::Span<DescriptorSetLayout* const> set_layouts,
- size_t push_constants);
- ~MetalPipelineArgumentBufferLayout() override;
-
- absl::Span<MetalArgumentBufferLayout* const> set_layouts() const {
- return set_layouts_;
- }
-
- private:
- absl::InlinedVector<MetalArgumentBufferLayout*, 2> set_layouts_;
- size_t push_constants_;
-};
-
-class MetalArgumentBuffer final : public DescriptorSet {
- public:
- MetalArgumentBuffer(MetalArgumentBufferLayout* layout,
- absl::Span<const Binding> resources);
- ~MetalArgumentBuffer() override;
-
- private:
- MetalArgumentBufferLayout* layout_;
- absl::InlinedVector<Binding, 8> resources_;
-};
-
-} // namespace metal
-} // namespace hal
-} // namespace iree
-
-#endif // IREE_HAL_METAL_METAL_PIPELINE_ARGUMENT_BUFFER_H_
diff --git a/iree/hal/metal/metal_pipeline_cache.h b/iree/hal/metal/metal_pipeline_cache.h
deleted file mode 100644
index 7b74909..0000000
--- a/iree/hal/metal/metal_pipeline_cache.h
+++ /dev/null
@@ -1,47 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#ifndef IREE_HAL_METAL_METAL_PIPELINE_CACHE_H_
-#define IREE_HAL_METAL_METAL_PIPELINE_CACHE_H_
-
-#import <Metal/Metal.h>
-
-#include "iree/hal/cc/executable_cache.h"
-
-namespace iree {
-namespace hal {
-namespace metal {
-
-// An ExecutableCache implementation for Metal.
-class MetalPipelineCache final : public ExecutableCache {
- public:
- explicit MetalPipelineCache(id<MTLDevice> device);
- ~MetalPipelineCache() override;
-
- bool CanPrepareFormat(iree_hal_executable_format_t format) const override;
-
- StatusOr<ref_ptr<Executable>> PrepareExecutable(
- ExecutableLayout* executable_layout,
- iree_hal_executable_caching_mode_t mode,
- iree_const_byte_span_t executable_data) override;
-
- private:
- id<MTLDevice> metal_device_;
-};
-
-} // namespace metal
-} // namespace hal
-} // namespace iree
-
-#endif // IREE_HAL_METAL_METAL_PIPELINE_CACHE_H_
diff --git a/iree/hal/metal/metal_pipeline_cache.mm b/iree/hal/metal/metal_pipeline_cache.mm
deleted file mode 100644
index ec328e4..0000000
--- a/iree/hal/metal/metal_pipeline_cache.mm
+++ /dev/null
@@ -1,50 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#include "iree/hal/metal/metal_pipeline_cache.h"
-
-#include "iree/base/status.h"
-#include "iree/base/tracing.h"
-#include "iree/hal/api.h"
-#include "iree/hal/metal/metal_kernel_library.h"
-
-namespace iree {
-namespace hal {
-namespace metal {
-
-static const iree_hal_executable_format_t kExecutableFormatMetal =
- iree_hal_make_executable_format("MTLE");
-
-MetalPipelineCache::MetalPipelineCache(id<MTLDevice> device) : metal_device_([device retain]) {}
-
-MetalPipelineCache::~MetalPipelineCache() { [metal_device_ release]; }
-
-bool MetalPipelineCache::CanPrepareFormat(iree_hal_executable_format_t format) const {
- return format == kExecutableFormatMetal;
-}
-
-StatusOr<ref_ptr<Executable>> MetalPipelineCache::PrepareExecutable(
- ExecutableLayout* executable_layout, iree_hal_executable_caching_mode_t mode,
- iree_const_byte_span_t executable_data) {
- IREE_TRACE_SCOPE0("MetalPipelineCache::PrepareExecutable");
-
- // Create the Metal library (which may itself own many pipeline states).
- IREE_ASSIGN_OR_RETURN(auto executable, MetalKernelLibrary::Create(metal_device_, mode, executable_data));
-
- return executable;
-}
-
-} // namespace metal
-} // namespace hal
-} // namespace iree
diff --git a/iree/hal/metal/metal_shared_event.h b/iree/hal/metal/metal_shared_event.h
deleted file mode 100644
index 35c4926..0000000
--- a/iree/hal/metal/metal_shared_event.h
+++ /dev/null
@@ -1,68 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#ifndef IREE_HAL_METAL_METAL_SHARED_EVENT_H_
-#define IREE_HAL_METAL_METAL_SHARED_EVENT_H_
-
-#import <Metal/Metal.h>
-
-#include "absl/synchronization/mutex.h"
-#include "iree/hal/cc/semaphore.h"
-
-namespace iree {
-namespace hal {
-namespace metal {
-
-// A semaphore implementation for Metal that directly wraps a MTLSharedEvent.
-class MetalSharedEvent final : public Semaphore {
- public:
- // Creates a MetalSharedEvent with the given |initial_value|.
- static StatusOr<ref_ptr<Semaphore>> Create(
- id<MTLDevice> device, MTLSharedEventListener* event_listener,
- uint64_t initial_value);
-
- ~MetalSharedEvent() override;
-
- id<MTLSharedEvent> handle() const { return metal_handle_; }
-
- StatusOr<uint64_t> Query() override;
-
- Status Signal(uint64_t value) override;
-
- void Fail(Status status) override;
-
- Status Wait(uint64_t value, Time deadline_ns) override;
-
- private:
- MetalSharedEvent(id<MTLDevice> device, MTLSharedEventListener* event_listener,
- uint64_t initial_value);
-
- id<MTLSharedEvent> metal_handle_;
-
- // An event listener for waiting and signaling. Its lifetime is managed by
- // the parent device.
- MTLSharedEventListener* event_listener_;
-
- // NOTE: the MTLSharedEvent is the source of truth. We only need to access
- // this status (and thus take the lock) when we want to either signal failure
- // or query the status in the case of the semaphore being set to UINT64_MAX.
- mutable absl::Mutex status_mutex_;
- Status status_ ABSL_GUARDED_BY(status_mutex_);
-};
-
-} // namespace metal
-} // namespace hal
-} // namespace iree
-
-#endif // IREE_HAL_METAL_METAL_SHARED_EVENT_H_
diff --git a/iree/hal/metal/metal_shared_event.mm b/iree/hal/metal/metal_shared_event.mm
deleted file mode 100644
index 325c30a..0000000
--- a/iree/hal/metal/metal_shared_event.mm
+++ /dev/null
@@ -1,108 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#include "iree/hal/metal/metal_shared_event.h"
-
-#include "iree/base/status.h"
-#include "iree/base/tracing.h"
-#include "iree/hal/metal/dispatch_time_util.h"
-
-namespace iree {
-namespace hal {
-namespace metal {
-
-// static
-StatusOr<ref_ptr<Semaphore>> MetalSharedEvent::Create(id<MTLDevice> device,
- MTLSharedEventListener* event_listener,
- uint64_t initial_value) {
- return assign_ref(new MetalSharedEvent(device, event_listener, initial_value));
-}
-
-MetalSharedEvent::MetalSharedEvent(id<MTLDevice> device, MTLSharedEventListener* event_listener,
- uint64_t initial_value)
- : metal_handle_([device newSharedEvent]), event_listener_(event_listener) {
- IREE_TRACE_SCOPE0("MetalSharedEvent::ctor");
- metal_handle_.signaledValue = initial_value;
-}
-
-MetalSharedEvent::~MetalSharedEvent() {
- IREE_TRACE_SCOPE0("MetalSharedEvent::dtor");
- [metal_handle_ release];
-}
-
-StatusOr<uint64_t> MetalSharedEvent::Query() {
- IREE_TRACE_SCOPE0("MetalSharedEvent::Query");
- uint64_t value = metal_handle_.signaledValue;
- if (value == UINT64_MAX) {
- absl::MutexLock lock(&status_mutex_);
- return status_;
- }
- return value;
-}
-
-Status MetalSharedEvent::Signal(uint64_t value) {
- IREE_TRACE_SCOPE0("MetalSharedEvent::Signal");
- metal_handle_.signaledValue = value;
- return OkStatus();
-}
-
-void MetalSharedEvent::Fail(Status status) {
- IREE_TRACE_SCOPE0("MetalSharedEvent::Fail");
- absl::MutexLock lock(&status_mutex_);
- status_ = std::move(status);
- metal_handle_.signaledValue = UINT64_MAX;
-}
-
-Status MetalSharedEvent::Wait(uint64_t value, Time deadline_ns) {
- IREE_TRACE_SCOPE0("MetalSharedEvent::Wait");
-
- Duration duration_ns = DeadlineToRelativeTimeoutNanos(deadline_ns);
- dispatch_time_t timeout = DurationToDispatchTime(duration_ns);
-
- // Quick path for impatient waiting to avoid all the overhead of dispatch queues and semaphores.
- if (duration_ns == ZeroDuration()) {
- if (metal_handle_.signaledValue < value) {
- return DeadlineExceededErrorBuilder(IREE_LOC) << "Deadline exceeded waiting for semaphores";
- }
- return OkStatus();
- }
-
- // Theoretically we don't really need to mark the semaphore handle as __block given that the
- // handle itself is not modified and there is only one block and it will copy the handle.
- // But marking it as __block serves as good documentation purpose.
- __block dispatch_semaphore_t work_done = dispatch_semaphore_create(0);
-
- // Use a listener to the MTLSharedEvent to notify us when the work is done on GPU by signaling a
- // semaphore. The signaling will happen in a new dispatch queue; the current thread will wait on
- // the semaphore.
- [metal_handle_ notifyListener:event_listener_
- atValue:value
- block:^(id<MTLSharedEvent>, uint64_t) {
- dispatch_semaphore_signal(work_done);
- }];
-
- long timed_out = dispatch_semaphore_wait(work_done, timeout);
-
- dispatch_release(work_done);
-
- if (timed_out) {
- return DeadlineExceededErrorBuilder(IREE_LOC)
- << "Deadline exceeded waiting for dispatch_semaphore_t";
- }
- return OkStatus();
-}
-
-} // namespace metal
-} // namespace hal
-} // namespace iree
diff --git a/iree/hal/metal/registration/BUILD.bazel b/iree/hal/metal/registration/BUILD.bazel
deleted file mode 100644
index 6a87e62..0000000
--- a/iree/hal/metal/registration/BUILD.bazel
+++ /dev/null
@@ -1,52 +0,0 @@
-# Copyright 2020 Google LLC
-#
-# Licensed under the Apache License, Version 2.0 (the "License");
-# you may not use this file except in compliance with the License.
-# You may obtain a copy of the License at
-#
-# https://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing, software
-# distributed under the License is distributed on an "AS IS" BASIS,
-# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-# See the License for the specific language governing permissions and
-# limitations under the License.
-
-load("//iree:build_defs.oss.bzl", "iree_cmake_extra_content")
-
-package(
- default_visibility = ["//visibility:public"],
- features = ["layering_check"],
- licenses = ["notice"], # Apache 2.0
-)
-
-iree_cmake_extra_content(
- content = """
-if(${IREE_HAL_DRIVER_METAL})
-""",
- inline = True,
-)
-
-cc_library(
- name = "registration",
- srcs = ["driver_module.cc"],
- hdrs = ["driver_module.h"],
- defines = [
- "IREE_HAL_HAVE_METAL_DRIVER_MODULE=1",
- ],
- deps = [
- "//iree/base:flags",
- "//iree/base:status",
- "//iree/base:tracing",
- "//iree/hal:api",
- "//iree/hal/metal",
- "@com_google_absl//absl/flags:flag",
- ],
-)
-
-iree_cmake_extra_content(
- content = """
-endif()
-""",
- inline = True,
-)
diff --git a/iree/hal/metal/registration/CMakeLists.txt b/iree/hal/metal/registration/CMakeLists.txt
deleted file mode 100644
index 90033d6..0000000
--- a/iree/hal/metal/registration/CMakeLists.txt
+++ /dev/null
@@ -1,38 +0,0 @@
-# Copyright 2020 Google LLC
-#
-# Licensed under the Apache License, Version 2.0 (the "License");
-# you may not use this file except in compliance with the License.
-# You may obtain a copy of the License at
-#
-# https://www.apache.org/licenses/LICENSE-2.0
-#
-# Unless required by applicable law or agreed to in writing, software
-# distributed under the License is distributed on an "AS IS" BASIS,
-# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-# See the License for the specific language governing permissions and
-# limitations under the License.
-
-iree_add_all_subdirs()
-
-if(${IREE_HAL_DRIVER_METAL})
-
-iree_cc_library(
- NAME
- registration
- HDRS
- "driver_module.h"
- SRCS
- "driver_module.cc"
- DEPS
- absl::flags
- iree::base::flags
- iree::base::status
- iree::base::tracing
- iree::hal::api
- iree::hal::metal
- DEFINES
- "IREE_HAL_HAVE_METAL_DRIVER_MODULE=1"
- PUBLIC
-)
-
-endif()
diff --git a/iree/hal/metal/registration/driver_module.cc b/iree/hal/metal/registration/driver_module.cc
deleted file mode 100644
index 1a9da7d..0000000
--- a/iree/hal/metal/registration/driver_module.cc
+++ /dev/null
@@ -1,70 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#include "iree/hal/metal/registration/driver_module.h"
-
-#include <inttypes.h>
-
-#include "absl/flags/flag.h"
-#include "iree/base/flags.h"
-#include "iree/hal/metal/metal_driver.h"
-
-ABSL_FLAG(bool, metal_capture, false, "Enables capturing Metal commands.");
-ABSL_FLAG(
- std::string, metal_capture_to_file, "",
- "Full path to store the GPU trace file (empty means capture to Xcode)");
-
-#define IREE_HAL_METAL_DRIVER_ID 0x4D544C31u // MTL1
-
-static iree_status_t iree_hal_metal_driver_factory_enumerate(
- void* self, const iree_hal_driver_info_t** out_driver_infos,
- iree_host_size_t* out_driver_info_count) {
- // NOTE: we could query supported metal versions or featuresets here.
- static const iree_hal_driver_info_t driver_infos[1] = {{
- /*driver_id=*/IREE_HAL_METAL_DRIVER_ID,
- /*driver_name=*/iree_make_cstring_view("metal"),
- /*full_name=*/iree_make_cstring_view("Apple Metal GPU"),
- }};
- *out_driver_info_count = IREE_ARRAYSIZE(driver_infos);
- *out_driver_infos = driver_infos;
- return iree_ok_status();
-}
-
-static iree_status_t iree_hal_metal_driver_factory_try_create(
- void* self, iree_hal_driver_id_t driver_id, iree_allocator_t allocator,
- iree_hal_driver_t** out_driver) {
- if (driver_id != IREE_HAL_METAL_DRIVER_ID) {
- return iree_make_status(IREE_STATUS_UNAVAILABLE,
- "no driver with ID %016" PRIu64
- " is provided by this factory",
- driver_id);
- }
- iree::hal::metal::MetalDriverOptions options;
- options.enable_capture = absl::GetFlag(FLAGS_metal_capture);
- options.capture_file = absl::GetFlag(FLAGS_metal_capture_to_file);
- IREE_ASSIGN_OR_RETURN(auto driver,
- iree::hal::metal::MetalDriver::Create(options));
- *out_driver = reinterpret_cast<iree_hal_driver_t*>(driver.release());
- return iree_ok_status();
-}
-
-IREE_API_EXPORT iree_status_t IREE_API_CALL
-iree_hal_metal_driver_module_register(iree_hal_driver_registry_t* registry) {
- static const iree_hal_driver_factory_t factory = {
- /*self=*/NULL,
- iree_hal_metal_driver_factory_enumerate,
- iree_hal_metal_driver_factory_try_create,
- };
- return iree_hal_driver_registry_register_factory(registry, &factory);
-}
diff --git a/iree/hal/metal/registration/driver_module.h b/iree/hal/metal/registration/driver_module.h
deleted file mode 100644
index edb6c05..0000000
--- a/iree/hal/metal/registration/driver_module.h
+++ /dev/null
@@ -1,31 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#ifndef IREE_HAL_METAL_REGISTRATION_DRIVER_MODULE_H_
-#define IREE_HAL_METAL_REGISTRATION_DRIVER_MODULE_H_
-
-#include "iree/hal/api.h"
-
-#ifdef __cplusplus
-extern "C" {
-#endif // __cplusplus
-
-IREE_API_EXPORT iree_status_t IREE_API_CALL
-iree_hal_metal_driver_module_register(iree_hal_driver_registry_t* registry);
-
-#ifdef __cplusplus
-} // extern "C"
-#endif // __cplusplus
-
-#endif // IREE_HAL_METAL_REGISTRATION_DRIVER_MODULE_H_
diff --git a/iree/test/e2e/xla_ops/BUILD b/iree/test/e2e/xla_ops/BUILD
index a4abcb5..2693834 100644
--- a/iree/test/e2e/xla_ops/BUILD
+++ b/iree/test/e2e/xla_ops/BUILD
@@ -34,45 +34,6 @@
)
iree_check_single_backend_test_suite(
- name = "check_metal-spirv_metal",
- srcs = [
- "abs.mlir",
- "add.mlir",
- "broadcast.mlir",
- "broadcast_add.mlir",
- "broadcast_in_dim.mlir",
- "clamp.mlir",
- "compare.mlir",
- "constant.mlir",
- "convert.mlir",
- "cosine.mlir",
- "divide.mlir",
- "exponential.mlir",
- "gather.mlir",
- "log.mlir",
- "log_plus_one.mlir",
- "maximum.mlir",
- "minimum.mlir",
- "multiply.mlir",
- "negate.mlir",
- "remainder.mlir",
- "reshape.mlir",
- "rsqrt.mlir",
- "select.mlir",
- "sine.mlir",
- "slice.mlir",
- "sqrt.mlir",
- "subtract.mlir",
- "tanh.mlir",
- "torch_index_select.mlir",
- "transpose.mlir",
- "while.mlir",
- ],
- driver = "metal",
- target_backend = "metal-spirv",
-)
-
-iree_check_single_backend_test_suite(
name = "check_vulkan-spirv_vulkan",
srcs = [
"abs.mlir",
@@ -177,7 +138,6 @@
name = "check",
tests = [
":check_dylib-llvm-aot_dylib",
- ":check_metal-spirv_metal",
":check_vmla_vmla",
":check_vulkan-spirv_vulkan",
],
diff --git a/iree/test/e2e/xla_ops/CMakeLists.txt b/iree/test/e2e/xla_ops/CMakeLists.txt
index 0c8978d..65c4d93 100644
--- a/iree/test/e2e/xla_ops/CMakeLists.txt
+++ b/iree/test/e2e/xla_ops/CMakeLists.txt
@@ -28,47 +28,6 @@
iree_check_single_backend_test_suite(
NAME
- check_metal-spirv_metal
- SRCS
- "abs.mlir"
- "add.mlir"
- "broadcast.mlir"
- "broadcast_add.mlir"
- "broadcast_in_dim.mlir"
- "clamp.mlir"
- "compare.mlir"
- "constant.mlir"
- "convert.mlir"
- "cosine.mlir"
- "divide.mlir"
- "exponential.mlir"
- "gather.mlir"
- "log.mlir"
- "log_plus_one.mlir"
- "maximum.mlir"
- "minimum.mlir"
- "multiply.mlir"
- "negate.mlir"
- "remainder.mlir"
- "reshape.mlir"
- "rsqrt.mlir"
- "select.mlir"
- "sine.mlir"
- "slice.mlir"
- "sqrt.mlir"
- "subtract.mlir"
- "tanh.mlir"
- "torch_index_select.mlir"
- "transpose.mlir"
- "while.mlir"
- TARGET_BACKEND
- "metal-spirv"
- DRIVER
- "metal"
-)
-
-iree_check_single_backend_test_suite(
- NAME
check_vulkan-spirv_vulkan
SRCS
"abs.mlir"
diff --git a/iree/test/e2e/xla_ops/partial/BUILD b/iree/test/e2e/xla_ops/partial/BUILD
index 9adea0c..3228fcd 100644
--- a/iree/test/e2e/xla_ops/partial/BUILD
+++ b/iree/test/e2e/xla_ops/partial/BUILD
@@ -39,14 +39,6 @@
)
iree_check_single_backend_test_suite(
- name = "check_metal-spirv_metal",
- srcs = [
- ],
- driver = "metal",
- target_backend = "metal-spirv",
-)
-
-iree_check_single_backend_test_suite(
name = "check_vulkan-spirv_vulkan",
srcs = [
],
@@ -66,7 +58,6 @@
name = "check",
tests = [
":check_dylib-llvm-aot_dylib",
- ":check_metal-spirv_metal",
":check_vmla_vmla",
":check_vulkan-spirv_vulkan",
],
diff --git a/iree/test/e2e/xla_ops/partial/CMakeLists.txt b/iree/test/e2e/xla_ops/partial/CMakeLists.txt
index 7cc0f00..43b2bba 100644
--- a/iree/test/e2e/xla_ops/partial/CMakeLists.txt
+++ b/iree/test/e2e/xla_ops/partial/CMakeLists.txt
@@ -28,15 +28,6 @@
iree_check_single_backend_test_suite(
NAME
- check_metal-spirv_metal
- TARGET_BACKEND
- "metal-spirv"
- DRIVER
- "metal"
-)
-
-iree_check_single_backend_test_suite(
- NAME
check_vulkan-spirv_vulkan
TARGET_BACKEND
"vulkan-spirv"
