[metal] Drop IREE HAL event via Metal fence implementation
IREE's HAL event is more flexible than what the Metal fence
supports, so we cannot just directly map there. We'd need
to find better ways to support. Though HAL even is mostly
unused right now so it's fine to be unimplemented.
diff --git a/experimental/metal/CMakeLists.txt b/experimental/metal/CMakeLists.txt
index 5e047ce..0b0fc22 100644
--- a/experimental/metal/CMakeLists.txt
+++ b/experimental/metal/CMakeLists.txt
@@ -27,8 +27,6 @@
"metal_buffer.m"
"metal_device.m"
"metal_driver.m"
- "metal_fence.h"
- "metal_fence.m"
"metal_kernel_library.h"
"metal_kernel_library.m"
"metal_shared_event.h"
diff --git a/experimental/metal/cts/CMakeLists.txt b/experimental/metal/cts/CMakeLists.txt
index 45c9c37..dedeb9f 100644
--- a/experimental/metal/cts/CMakeLists.txt
+++ b/experimental/metal/cts/CMakeLists.txt
@@ -17,5 +17,8 @@
"\"MTLE\""
DEPS
iree::experimental::metal::registration
+ EXCLUDED_TESTS
+ # HAL event is unimplemented for Metal right now.
+ "event"
)
diff --git a/experimental/metal/direct_command_buffer.m b/experimental/metal/direct_command_buffer.m
index 3a3c202..956f369 100644
--- a/experimental/metal/direct_command_buffer.m
+++ b/experimental/metal/direct_command_buffer.m
@@ -11,7 +11,6 @@
#include "experimental/metal/builtin_executables.h"
#include "experimental/metal/metal_buffer.h"
#include "experimental/metal/metal_device.h"
-#include "experimental/metal/metal_fence.h"
#include "experimental/metal/metal_kernel_library.h"
#include "experimental/metal/pipeline_layout.h"
#include "iree/base/api.h"
@@ -311,30 +310,13 @@
static iree_status_t iree_hal_metal_command_buffer_signal_event(
iree_hal_command_buffer_t* base_command_buffer, iree_hal_event_t* event,
iree_hal_execution_stage_t source_stage_mask) {
- iree_hal_metal_command_buffer_t* command_buffer =
- iree_hal_metal_command_buffer_cast(base_command_buffer);
- id<MTLFence> fence = iree_hal_metal_fence_handle(event);
-
- // In Metal compute pipelines, fences are more course grained than Vulkan--we only have one
- // execution stage instead of multiple stages in Vulkan.
- if (command_buffer->blit_encoder != nil) {
- [command_buffer->blit_encoder updateFence:fence];
- } else {
- id<MTLComputeCommandEncoder> compute_encoder =
- iree_hal_metal_get_or_begin_compute_encoder(command_buffer);
- [compute_encoder updateFence:fence];
- }
-
- return iree_ok_status();
+ return iree_make_status(IREE_STATUS_UNIMPLEMENTED, "event not yet supported");
}
static iree_status_t iree_hal_metal_command_buffer_reset_event(
iree_hal_command_buffer_t* base_command_buffer, iree_hal_event_t* event,
iree_hal_execution_stage_t source_stage_mask) {
- // In Metal, MTLFence does not support reset. We just create a new fence every time IREE event
- // reset API is called. This assumes IREE fences are not waited anymore when reset is called,
- // which should be true for proper IREE event usage.
- return iree_hal_metal_fence_recreate(event);
+ return iree_make_status(IREE_STATUS_UNIMPLEMENTED, "event not yet supported");
}
static iree_status_t iree_hal_metal_command_buffer_wait_events(
@@ -343,27 +325,7 @@
iree_hal_execution_stage_t target_stage_mask, iree_host_size_t memory_barrier_count,
const iree_hal_memory_barrier_t* memory_barriers, iree_host_size_t buffer_barrier_count,
const iree_hal_buffer_barrier_t* buffer_barriers) {
- iree_hal_metal_command_buffer_t* command_buffer =
- iree_hal_metal_command_buffer_cast(base_command_buffer);
-
- // In Metal compute pipelines, fences are more course grained than Vulkan--we only have one
- // execution stage instead of multiple stages in Vulkan, and we don't have separate memory
- // barriers to control.
- if (command_buffer->blit_encoder != nil) {
- for (iree_host_size_t i = 0; i < event_count; ++i) {
- id<MTLFence> fence = iree_hal_metal_fence_handle(events[i]);
- [command_buffer->blit_encoder waitForFence:fence];
- }
- } else {
- id<MTLComputeCommandEncoder> compute_encoder =
- iree_hal_metal_get_or_begin_compute_encoder(command_buffer);
- for (iree_host_size_t i = 0; i < event_count; ++i) {
- id<MTLFence> fence = iree_hal_metal_fence_handle(events[i]);
- [compute_encoder waitForFence:fence];
- }
- }
-
- return iree_ok_status();
+ return iree_make_status(IREE_STATUS_UNIMPLEMENTED, "event not yet supported");
}
static iree_status_t iree_hal_metal_command_buffer_discard_buffer(
diff --git a/experimental/metal/metal_device.m b/experimental/metal/metal_device.m
index 67a07fd..10b166c 100644
--- a/experimental/metal/metal_device.m
+++ b/experimental/metal/metal_device.m
@@ -10,7 +10,6 @@
#include "experimental/metal/builtin_executables.h"
#include "experimental/metal/direct_allocator.h"
#include "experimental/metal/direct_command_buffer.h"
-#include "experimental/metal/metal_fence.h"
#include "experimental/metal/metal_shared_event.h"
#include "experimental/metal/nop_executable_cache.h"
#include "experimental/metal/pipeline_layout.h"
@@ -237,8 +236,7 @@
static iree_status_t iree_hal_metal_device_create_event(iree_hal_device_t* base_device,
iree_hal_event_t** out_event) {
- iree_hal_metal_device_t* device = iree_hal_metal_device_cast(base_device);
- return iree_hal_metal_fence_create(device->device, device->host_allocator, out_event);
+ return iree_make_status(IREE_STATUS_UNIMPLEMENTED, "event not yet supported");
}
static iree_status_t iree_hal_metal_device_create_executable_cache(
diff --git a/experimental/metal/metal_fence.h b/experimental/metal/metal_fence.h
deleted file mode 100644
index 9a7c72b..0000000
--- a/experimental/metal/metal_fence.h
+++ /dev/null
@@ -1,38 +0,0 @@
-// Copyright 2023 The IREE Authors
-//
-// Licensed under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-
-#ifndef IREE_EXPERIMENTAL_METAL_METAL_FENCE_H_
-#define IREE_EXPERIMENTAL_METAL_METAL_FENCE_H_
-
-#import <Metal/Metal.h>
-
-#include "iree/base/api.h"
-#include "iree/hal/api.h"
-
-#ifdef __cplusplus
-extern "C" {
-#endif // __cplusplus
-
-// Creates a Metal fence to implement an IREE event.
-//
-// An IREE event defines synchronization scopes within the same command buffer.
-// It maps to MTLFence in Metal.
-iree_status_t iree_hal_metal_fence_create(id<MTLDevice> device,
- iree_allocator_t host_allocator,
- iree_hal_event_t** out_event);
-
-// Destroys the current Metal fence behind the given |event| and recreate a new
-// one. This is meant to support the IREE event reset API.
-iree_status_t iree_hal_metal_fence_recreate(iree_hal_event_t* event);
-
-// Returns the underlying Metal fence handle for the given |base_event|.
-id<MTLFence> iree_hal_metal_fence_handle(const iree_hal_event_t* base_event);
-
-#ifdef __cplusplus
-} // extern "C"
-#endif // __cplusplus
-
-#endif // IREE_EXPERIMENTAL_METAL_METAL_FENCE_H_
diff --git a/experimental/metal/metal_fence.m b/experimental/metal/metal_fence.m
deleted file mode 100644
index 72bf8a8..0000000
--- a/experimental/metal/metal_fence.m
+++ /dev/null
@@ -1,84 +0,0 @@
-// Copyright 2023 The IREE Authors
-//
-// Licensed under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-
-#include "experimental/metal/metal_fence.h"
-
-#import <Metal/Metal.h>
-
-#include "iree/base/api.h"
-#include "iree/base/tracing.h"
-
-typedef struct iree_hal_metal_fence_t {
- // Abstract resource used for injecting reference counting and vtable; must be at offset 0.
- iree_hal_resource_t resource;
-
- id<MTLFence> fence;
-
- iree_allocator_t host_allocator;
-} iree_hal_metal_fence_t;
-
-static const iree_hal_event_vtable_t iree_hal_metal_fence_vtable;
-
-static iree_hal_metal_fence_t* iree_hal_metal_fence_cast(iree_hal_event_t* base_value) {
- IREE_HAL_ASSERT_TYPE(base_value, &iree_hal_metal_fence_vtable);
- return (iree_hal_metal_fence_t*)base_value;
-}
-
-static const iree_hal_metal_fence_t* iree_hal_metal_fence_const_cast(
- const iree_hal_event_t* base_value) {
- IREE_HAL_ASSERT_TYPE(base_value, &iree_hal_metal_fence_vtable);
- return (const iree_hal_metal_fence_t*)base_value;
-}
-
-id<MTLFence> iree_hal_metal_fence_handle(const iree_hal_event_t* base_event) {
- const iree_hal_metal_fence_t* fence = iree_hal_metal_fence_const_cast(base_event);
- return fence->fence;
-}
-
-iree_status_t iree_hal_metal_fence_create(id<MTLDevice> device, iree_allocator_t host_allocator,
- iree_hal_event_t** out_event) {
- IREE_ASSERT_ARGUMENT(out_event);
- IREE_TRACE_ZONE_BEGIN(z0);
-
- *out_event = NULL;
- iree_hal_metal_fence_t* event = NULL;
- iree_status_t status = iree_allocator_malloc(host_allocator, sizeof(*event), (void**)&event);
- if (iree_status_is_ok(status)) {
- iree_hal_resource_initialize(&iree_hal_metal_fence_vtable, &event->resource);
- event->fence = [device newFence]; // +1
- event->host_allocator = host_allocator;
- *out_event = (iree_hal_event_t*)event;
- }
-
- IREE_TRACE_ZONE_END(z0);
- return status;
-}
-
-iree_status_t iree_hal_metal_fence_recreate(iree_hal_event_t* event) {
- iree_hal_metal_fence_t* fence = iree_hal_metal_fence_cast(event);
- IREE_TRACE_ZONE_BEGIN(z0);
-
- id<MTLDevice> device = fence->fence.device;
- [fence->fence release];
- fence->fence = [device newFence];
-
- IREE_TRACE_ZONE_END(z0);
- return iree_ok_status();
-}
-
-static void iree_hal_metal_fence_destroy(iree_hal_event_t* base_event) {
- iree_hal_metal_fence_t* fence = iree_hal_metal_fence_cast(base_event);
- IREE_TRACE_ZONE_BEGIN(z0);
-
- [fence->fence release]; // -1
- iree_allocator_free(fence->host_allocator, base_event);
-
- IREE_TRACE_ZONE_END(z0);
-}
-
-static const iree_hal_event_vtable_t iree_hal_metal_fence_vtable = {
- .destroy = iree_hal_metal_fence_destroy,
-};
