commit ffb40b14e9ffd30ead8848a9c09b1b99ed313809
author Lei Zhang <antiagainst@google.com> Thu Apr 27 08:18:51 2023 -0700
committer Lei Zhang <antiagainst@google.com> Tue Jun 13 21:17:32 2023 -0700
tree c24a191773d72eeb9d6d072e44f154cddc32f4c5
parent 30d52cda78c63604fa9db0461afdadcd57ada5f8

[metal] Signal and wait MTLEvent for execution only barriers

There is no direct corresponding APIs for execution only barrier
in Metal. We just signal and wait on the same value of a MTLEvent.

experimental/metal/direct_command_buffer.m

diff --git a/experimental/metal/direct_command_buffer.m b/experimental/metal/direct_command_buffer.m
index 1eaf951..62fe3d1 100644
--- a/experimental/metal/direct_command_buffer.m
+++ b/experimental/metal/direct_command_buffer.m
@@ -291,6 +291,19 @@
 
   iree_hal_metal_command_buffer_t* command_buffer =
       iree_hal_metal_command_buffer_cast(base_command_buffer);
+
+  if (memory_barrier_count == 0 && buffer_barrier_count == 0) {
+    // There is no direct corresponding APIs for execution only barrier in Metal. We just signal and
+    // wait on the same value of a MTLEvent here.
+    iree_hal_metal_end_blit_encoder(command_buffer);
+    iree_hal_metal_end_compute_encoder(command_buffer);
+    id<MTLCommandBuffer> metal_handle = command_buffer->command_buffer;
+    uint64_t event_value = command_buffer->next_encoder_event_value++;
+    [metal_handle encodeSignalEvent:command_buffer->encoder_event value:event_value];
+    [metal_handle encodeWaitForEvent:command_buffer->encoder_event value:event_value];
+    return iree_ok_status();
+  }
+
   id<MTLComputeCommandEncoder> encoder =
       iree_hal_metal_get_or_begin_compute_encoder(command_buffer);