Metal: Implement begin/end access synchronization with MTLSharedEvent

Bug: b/252731382
Change-Id: Ie2bf978c10dcb7b2c03a2c7ff81ddd8b9b77ac20
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/106760
Reviewed-by: Shrek Shao <shrekshao@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Austin Eng <enga@chromium.org>

include/dawn/native/MetalBackend.h

@@ -15,6 +15,8 @@
 #ifndef INCLUDE_DAWN_NATIVE_METALBACKEND_H_
 #define INCLUDE_DAWN_NATIVE_METALBACKEND_H_
 
+#include <vector>
+
 #include "dawn/dawn_wsi.h"
 #include "dawn/native/DawnNative.h"
 
@@ -38,19 +40,47 @@ struct DAWN_NATIVE_EXPORT AdapterDiscoveryOptions : public AdapterDiscoveryOptio
     AdapterDiscoveryOptions();
 };
 
+struct DAWN_NATIVE_EXPORT ExternalImageMTLSharedEventDescriptor {
+    // Shared event handle `id<MTLSharedEvent>`.
+    // This never passes ownership to the callee (when used as an input
+    // parameter) or to the caller (when used as a return value or output parameter).
+#ifdef __OBJC__
+    id<MTLSharedEvent> sharedEvent = nil;
+    static_assert(sizeof(id<MTLSharedEvent>) == sizeof(void*));
+    static_assert(alignof(id<MTLSharedEvent>) == alignof(void*));
+#else
+    void* sharedEvent = nullptr;
+#endif
+
+    // The value that was previously signaled on this event and should be waited on.
+    uint64_t signaledValue = 0;
+};
+
 struct DAWN_NATIVE_EXPORT ExternalImageDescriptorIOSurface : ExternalImageDescriptor {
   public:
     ExternalImageDescriptorIOSurface();
+    ~ExternalImageDescriptorIOSurface();
 
     IOSurfaceRef ioSurface;
 
     // This has been deprecated.
     uint32_t plane;
+
+    // A list of events to wait on before accessing the texture.
+    std::vector<ExternalImageMTLSharedEventDescriptor> waitEvents;
+};
+
+struct DAWN_NATIVE_EXPORT ExternalImageIOSurfaceEndAccessDescriptor
+    : ExternalImageMTLSharedEventDescriptor {
+    bool isInitialized;
 };
 
 DAWN_NATIVE_EXPORT WGPUTexture WrapIOSurface(WGPUDevice device,
                                              const ExternalImageDescriptorIOSurface* descriptor);
 
+DAWN_NATIVE_EXPORT void IOSurfaceEndAccess(WGPUTexture texture,
+                                           ExternalImageIOSurfaceEndAccessDescriptor* descriptor);
+
 // When making Metal interop with other APIs, we need to be careful that QueueSubmit doesn't
 // mean that the operations will be visible to other APIs/Metal devices right away. macOS
 // does have a global queue of graphics operations, but the command buffers are inserted there

src/dawn/native/metal/CommandBufferMTL.mm

@@ -761,6 +761,10 @@ MaybeError CommandBuffer::FillCommands(CommandRecordingContext* commandContext)
             case Command::BeginComputePass: {
                 BeginComputePassCmd* cmd = mCommands.NextCommand<BeginComputePassCmd>();
 
+                for (TextureBase* texture :
+                     GetResourceUsages().computePasses[nextComputePassNumber].referencedTextures) {
+                    ToBackend(texture)->SynchronizeTextureBeforeUse(commandContext);
+                }
                 for (const SyncScopeResourceUsage& scope :
                      GetResourceUsages().computePasses[nextComputePassNumber].dispatchUsages) {
                     LazyClearSyncScope(scope, commandContext);
@@ -776,6 +780,20 @@ MaybeError CommandBuffer::FillCommands(CommandRecordingContext* commandContext)
             case Command::BeginRenderPass: {
                 BeginRenderPassCmd* cmd = mCommands.NextCommand<BeginRenderPassCmd>();
 
+                for (TextureBase* texture :
+                     this->GetResourceUsages().renderPasses[nextRenderPassNumber].textures) {
+                    ToBackend(texture)->SynchronizeTextureBeforeUse(commandContext);
+                }
+                for (ExternalTextureBase* externalTexture : this->GetResourceUsages()
+                                                                .renderPasses[nextRenderPassNumber]
+                                                                .externalTextures) {
+                    for (auto& view : externalTexture->GetTextureViews()) {
+                        if (view.Get()) {
+                            Texture* texture = ToBackend(view->GetTexture());
+                            texture->SynchronizeTextureBeforeUse(commandContext);
+                        }
+                    }
+                }
                 LazyClearSyncScope(GetResourceUsages().renderPasses[nextRenderPassNumber],
                                    commandContext);
                 commandContext->EndBlit();
@@ -831,6 +849,7 @@ MaybeError CommandBuffer::FillCommands(CommandRecordingContext* commandContext)
                 buffer->EnsureDataInitialized(commandContext);
                 EnsureDestinationTextureInitialized(commandContext, texture, dst, copySize);
 
+                texture->SynchronizeTextureBeforeUse(commandContext);
                 RecordCopyBufferToTexture(commandContext, buffer->GetMTLBuffer(), buffer->GetSize(),
                                           src.offset, src.bytesPerRow, src.rowsPerImage, texture,
                                           dst.mipLevel, dst.origin, dst.aspect, copySize);
@@ -852,6 +871,7 @@ MaybeError CommandBuffer::FillCommands(CommandRecordingContext* commandContext)
 
                 buffer->EnsureDataInitializedAsDestination(commandContext, copy);
 
+                texture->SynchronizeTextureBeforeUse(commandContext);
                 texture->EnsureSubresourceContentInitialized(
                     commandContext, GetSubresourcesAffectedByCopy(src, copySize));
 
@@ -941,6 +961,8 @@ MaybeError CommandBuffer::FillCommands(CommandRecordingContext* commandContext)
                 Texture* srcTexture = ToBackend(copy->source.texture.Get());
                 Texture* dstTexture = ToBackend(copy->destination.texture.Get());
 
+                srcTexture->SynchronizeTextureBeforeUse(commandContext);
+                dstTexture->SynchronizeTextureBeforeUse(commandContext);
                 srcTexture->EnsureSubresourceContentInitialized(
                     commandContext, GetSubresourcesAffectedByCopy(copy->source, copy->copySize));
                 EnsureDestinationTextureInitialized(commandContext, dstTexture, copy->destination,

src/dawn/native/metal/CommandRecordingContext.h

@@ -22,6 +22,11 @@
 
 namespace dawn::native::metal {
 
+struct MTLSharedEventAndSignalValue {
+    NSPRef<id> sharedEvent;
+    uint64_t signaledValue;
+};
+
 // This class wraps a MTLCommandBuffer and tracks which Metal encoder is open.
 // Only one encoder may be open at a time.
 class CommandRecordingContext : NonMovable {

src/dawn/native/metal/DeviceMTL.h

@@ -18,6 +18,7 @@
 #include <atomic>
 #include <memory>
 #include <mutex>
+#include <vector>
 
 #include "dawn/native/dawn_platform.h"
 
@@ -33,6 +34,7 @@
 namespace dawn::native::metal {
 
 struct KalmanInfo;
+struct ExternalImageMTLSharedEventDescriptor;
 
 class Device final : public DeviceBase {
   public:
@@ -53,8 +55,12 @@ class Device final : public DeviceBase {
         Device::SubmitMode submitMode = Device::SubmitMode::Normal);
     MaybeError SubmitPendingCommandBuffer();
 
-    Ref<Texture> CreateTextureWrappingIOSurface(const ExternalImageDescriptor* descriptor,
+    void ExportLastSignaledEvent(ExternalImageMTLSharedEventDescriptor* desc);
-                                                IOSurfaceRef ioSurface);
+
+    Ref<Texture> CreateTextureWrappingIOSurface(
+        const ExternalImageDescriptor* descriptor,
+        IOSurfaceRef ioSurface,
+        std::vector<MTLSharedEventAndSignalValue> waitEvents);
     void WaitForCommandsToBeScheduled();
 
     ResultOrError<std::unique_ptr<StagingBufferBase>> CreateStagingBuffer(size_t size) override;
@@ -134,6 +140,7 @@ class Device final : public DeviceBase {
     ResultOrError<ExecutionSerial> CheckAndUpdateCompletedSerials() override;
 
     NSPRef<id<MTLDevice>> mMtlDevice;
+    NSPRef<id> mMtlSharedEvent = nil;  // MTLSharedEvent not available until macOS 10.14+.
     NSPRef<id<MTLCommandQueue>> mCommandQueue;
 
     CommandRecordingContext mCommandContext;

src/dawn/native/metal/DeviceMTL.mm

@@ -146,6 +146,10 @@ MaybeError Device::Initialize(const DeviceDescriptor* descriptor) {
         return DAWN_INTERNAL_ERROR("Failed to allocate MTLCommandQueue.");
     }
 
+    if (@available(macOS 10.14, *)) {
+        mMtlSharedEvent.Acquire([*mMtlDevice newSharedEvent]);
+    }
+
     DAWN_TRY(mCommandContext.PrepareNextCommandBuffer(*mCommandQueue));
 
     if (HasFeature(Feature::TimestampQuery) &&
@@ -428,11 +432,21 @@ MaybeError Device::SubmitPendingCommandBuffer() {
 
     TRACE_EVENT_ASYNC_BEGIN0(GetPlatform(), GPUWork, "DeviceMTL::SubmitPendingCommandBuffer",
                              uint64_t(pendingSerial));
+    if (@available(macOS 10.14, *)) {
+        id rawEvent = *mMtlSharedEvent;
+        id<MTLSharedEvent> sharedEvent = static_cast<id<MTLSharedEvent>>(rawEvent);
+        [*pendingCommands encodeSignalEvent:sharedEvent value:static_cast<uint64_t>(pendingSerial)];
+    }
     [*pendingCommands commit];
 
     return mCommandContext.PrepareNextCommandBuffer(*mCommandQueue);
 }
 
+void Device::ExportLastSignaledEvent(ExternalImageMTLSharedEventDescriptor* desc) {
+    desc->sharedEvent = *mMtlSharedEvent;
+    desc->signaledValue = static_cast<uint64_t>(GetLastSubmittedCommandSerial());
+}
+
 ResultOrError<std::unique_ptr<StagingBufferBase>> Device::CreateStagingBuffer(size_t size) {
     std::unique_ptr<StagingBufferBase> stagingBuffer = std::make_unique<StagingBuffer>(size, this);
     DAWN_TRY(stagingBuffer->Initialize());
@@ -471,6 +485,7 @@ MaybeError Device::CopyFromStagingToTextureImpl(const StagingBufferBase* source,
                                                 TextureCopy* dst,
                                                 const Extent3D& copySizePixels) {
     Texture* texture = ToBackend(dst->texture.Get());
+    texture->SynchronizeTextureBeforeUse(GetPendingCommandContext());
     EnsureDestinationTextureInitialized(GetPendingCommandContext(DeviceBase::SubmitMode::Passive),
                                         texture, *dst, copySizePixels);
 
@@ -481,8 +496,10 @@ MaybeError Device::CopyFromStagingToTextureImpl(const StagingBufferBase* source,
     return {};
 }
 
-Ref<Texture> Device::CreateTextureWrappingIOSurface(const ExternalImageDescriptor* descriptor,
+Ref<Texture> Device::CreateTextureWrappingIOSurface(
-                                                    IOSurfaceRef ioSurface) {
+    const ExternalImageDescriptor* descriptor,
+    IOSurfaceRef ioSurface,
+    std::vector<MTLSharedEventAndSignalValue> waitEvents) {
     const TextureDescriptor* textureDescriptor = FromAPI(descriptor->cTextureDescriptor);
     if (ConsumedError(ValidateIsAlive())) {
         return nullptr;
@@ -495,7 +512,9 @@ Ref<Texture> Device::CreateTextureWrappingIOSurface(const ExternalImageDescripto
     }
 
     Ref<Texture> result;
-    if (ConsumedError(Texture::CreateFromIOSurface(this, descriptor, ioSurface), &result)) {
+    if (ConsumedError(
+            Texture::CreateFromIOSurface(this, descriptor, ioSurface, std::move(waitEvents)),
+            &result)) {
         return nullptr;
     }
     return result;

src/dawn/native/metal/MetalBackend.mm

@@ -17,6 +17,7 @@
 
 #include "dawn/native/MetalBackend.h"
 
+#include "dawn/native/metal/CommandRecordingContext.h"
 #include "dawn/native/metal/DeviceMTL.h"
 #include "dawn/native/metal/TextureMTL.h"
 
@@ -28,13 +29,26 @@ AdapterDiscoveryOptions::AdapterDiscoveryOptions()
 ExternalImageDescriptorIOSurface::ExternalImageDescriptorIOSurface()
     : ExternalImageDescriptor(ExternalImageType::IOSurface) {}
 
+ExternalImageDescriptorIOSurface::~ExternalImageDescriptorIOSurface() = default;
+
 WGPUTexture WrapIOSurface(WGPUDevice device, const ExternalImageDescriptorIOSurface* cDescriptor) {
     Device* backendDevice = ToBackend(FromAPI(device));
-    Ref<TextureBase> texture =
+    std::vector<MTLSharedEventAndSignalValue> waitEvents;
-        backendDevice->CreateTextureWrappingIOSurface(cDescriptor, cDescriptor->ioSurface);
+    for (const auto& waitEvent : cDescriptor->waitEvents) {
+        waitEvents.push_back(
+            {static_cast<id<MTLSharedEvent>>(waitEvent.sharedEvent), waitEvent.signaledValue});
+    }
+    Ref<TextureBase> texture = backendDevice->CreateTextureWrappingIOSurface(
+        cDescriptor, cDescriptor->ioSurface, std::move(waitEvents));
     return ToAPI(texture.Detach());
 }
 
+void IOSurfaceEndAccess(WGPUTexture cTexture,
+                        ExternalImageIOSurfaceEndAccessDescriptor* descriptor) {
+    Texture* texture = ToBackend(FromAPI(cTexture));
+    texture->IOSurfaceEndAccess(descriptor);
+}
+
 void WaitForCommandsToBeScheduled(WGPUDevice device) {
     ToBackend(FromAPI(device))->WaitForCommandsToBeScheduled();
 }

src/dawn/native/metal/TextureMTL.h

@@ -17,17 +17,20 @@
 
 #include <IOSurface/IOSurfaceRef.h>
 #import <Metal/Metal.h>
+#include <vector>
 
 #include "dawn/native/Texture.h"
 
 #include "dawn/common/CoreFoundationRef.h"
 #include "dawn/common/NSRef.h"
 #include "dawn/native/DawnNative.h"
+#include "dawn/native/MetalBackend.h"
 
 namespace dawn::native::metal {
 
 class CommandRecordingContext;
 class Device;
+struct MTLSharedEventAndSignalValue;
 
 MTLPixelFormat MetalPixelFormat(wgpu::TextureFormat format);
 MaybeError ValidateIOSurfaceCanBeWrapped(const DeviceBase* device,
@@ -40,7 +43,8 @@ class Texture final : public TextureBase {
     static ResultOrError<Ref<Texture>> CreateFromIOSurface(
         Device* device,
         const ExternalImageDescriptor* descriptor,
-        IOSurfaceRef ioSurface);
+        IOSurfaceRef ioSurface,
+        std::vector<MTLSharedEventAndSignalValue> waitEvents);
     static Ref<Texture> CreateWrapping(Device* device,
                                        const TextureDescriptor* descriptor,
                                        NSPRef<id<MTLTexture>> wrapped);
@@ -54,6 +58,9 @@ class Texture final : public TextureBase {
     void EnsureSubresourceContentInitialized(CommandRecordingContext* commandContext,
                                              const SubresourceRange& range);
 
+    void SynchronizeTextureBeforeUse(CommandRecordingContext* commandContext);
+    void IOSurfaceEndAccess(ExternalImageIOSurfaceEndAccessDescriptor* descriptor);
+
   private:
     using TextureBase::TextureBase;
     ~Texture() override;
@@ -63,7 +70,8 @@ class Texture final : public TextureBase {
     MaybeError InitializeAsInternalTexture(const TextureDescriptor* descriptor);
     MaybeError InitializeFromIOSurface(const ExternalImageDescriptor* descriptor,
                                        const TextureDescriptor* textureDescriptor,
-                                       IOSurfaceRef ioSurface);
+                                       IOSurfaceRef ioSurface,
+                                       std::vector<MTLSharedEventAndSignalValue> waitEvents);
     void InitializeAsWrapping(const TextureDescriptor* descriptor, NSPRef<id<MTLTexture>> wrapped);
 
     void DestroyImpl() override;
@@ -76,6 +84,7 @@ class Texture final : public TextureBase {
 
     MTLTextureUsage mMtlUsage;
     CFRef<IOSurfaceRef> mIOSurface = nullptr;
+    std::vector<MTLSharedEventAndSignalValue> mWaitEvents;
 };
 
 class TextureView final : public TextureViewBase {

src/dawn/native/metal/TextureMTL.mm

@@ -691,14 +691,17 @@ ResultOrError<Ref<Texture>> Texture::Create(Device* device, const TextureDescrip
 }
 
 // static
-ResultOrError<Ref<Texture>> Texture::CreateFromIOSurface(Device* device,
+ResultOrError<Ref<Texture>> Texture::CreateFromIOSurface(
+    Device* device,
     const ExternalImageDescriptor* descriptor,
-                                                         IOSurfaceRef ioSurface) {
+    IOSurfaceRef ioSurface,
+    std::vector<MTLSharedEventAndSignalValue> waitEvents) {
     const TextureDescriptor* textureDescriptor = FromAPI(descriptor->cTextureDescriptor);
 
     Ref<Texture> texture =
         AcquireRef(new Texture(device, textureDescriptor, TextureState::OwnedExternal));
-    DAWN_TRY(texture->InitializeFromIOSurface(descriptor, textureDescriptor, ioSurface));
+    DAWN_TRY(texture->InitializeFromIOSurface(descriptor, textureDescriptor, ioSurface,
+                                              std::move(waitEvents)));
     return texture;
 }
 
@@ -739,8 +742,10 @@ void Texture::InitializeAsWrapping(const TextureDescriptor* descriptor,
 
 MaybeError Texture::InitializeFromIOSurface(const ExternalImageDescriptor* descriptor,
                                             const TextureDescriptor* textureDescriptor,
-                                            IOSurfaceRef ioSurface) {
+                                            IOSurfaceRef ioSurface,
+                                            std::vector<MTLSharedEventAndSignalValue> waitEvents) {
     mIOSurface = ioSurface;
+    mWaitEvents = std::move(waitEvents);
 
     // Uses WGPUTexture which wraps multiplanar ioSurface needs to create
     // texture view explicitly. Wrap the ioSurface and delay to extract
@@ -763,6 +768,31 @@ MaybeError Texture::InitializeFromIOSurface(const ExternalImageDescriptor* descr
     return {};
 }
 
+void Texture::SynchronizeTextureBeforeUse(CommandRecordingContext* commandContext) {
+    if (@available(macOS 10.14, *)) {
+        if (!mWaitEvents.empty()) {
+            // There may be an open blit encoder from a copy command or writeBuffer.
+            // Wait events are only allowed if there is no encoder open.
+            commandContext->EndBlit();
+        }
+        auto commandBuffer = commandContext->GetCommands();
+        // Consume the wait events on the texture. They will be empty after this loop.
+        for (auto waitEvent : std::move(mWaitEvents)) {
+            id rawEvent = *waitEvent.sharedEvent;
+            id<MTLSharedEvent> sharedEvent = static_cast<id<MTLSharedEvent>>(rawEvent);
+            [commandBuffer encodeWaitForEvent:sharedEvent value:waitEvent.signaledValue];
+        }
+    }
+}
+
+void Texture::IOSurfaceEndAccess(ExternalImageIOSurfaceEndAccessDescriptor* descriptor) {
+    ASSERT(descriptor);
+    ToBackend(GetDevice())->ExportLastSignaledEvent(descriptor);
+    descriptor->isInitialized = IsSubresourceContentInitialized(GetAllSubresources());
+    // Destroy the texture as it should not longer be used after EndAccess.
+    Destroy();
+}
+
 Texture::Texture(DeviceBase* dev, const TextureDescriptor* desc, TextureState st)
     : TextureBase(dev, desc, st) {}
 

src/dawn/tests/end2end/IOSurfaceWrappingTests.cpp

@@ -449,6 +449,143 @@ TEST_P(IOSurfaceUsageTests, UninitializedTextureIsCleared) {
     // wrap ioSurface and ensure color is not visible when isInitialized set to false
     wgpu::Texture ioSurfaceTexture = WrapIOSurface(&textureDescriptor, ioSurface.get(), false);
     EXPECT_PIXEL_RGBA8_EQ(utils::RGBA8(0, 0, 0, 0), ioSurfaceTexture, 0, 0);
+
+    dawn::native::metal::ExternalImageIOSurfaceEndAccessDescriptor endAccessDesc;
+    dawn::native::metal::IOSurfaceEndAccess(ioSurfaceTexture.Get(), &endAccessDesc);
+    EXPECT_TRUE(endAccessDesc.isInitialized);
+}
+
+// Test that exporting a texture wrapping an IOSurface sets the isInitialized bit to
+// false if the contents are discard.
+TEST_P(IOSurfaceUsageTests, UninitializedOnEndAccess) {
+    DAWN_TEST_UNSUPPORTED_IF(UsesWire());
+
+    ScopedIOSurfaceRef ioSurface = CreateSinglePlaneIOSurface(1, 1, kCVPixelFormatType_32RGBA, 4);
+    uint32_t data = 0x04030201;
+
+    IOSurfaceLock(ioSurface.get(), 0, nullptr);
+    memcpy(IOSurfaceGetBaseAddress(ioSurface.get()), &data, sizeof(data));
+    IOSurfaceUnlock(ioSurface.get(), 0, nullptr);
+
+    wgpu::TextureDescriptor textureDescriptor;
+    textureDescriptor.dimension = wgpu::TextureDimension::e2D;
+    textureDescriptor.format = wgpu::TextureFormat::RGBA8Unorm;
+    textureDescriptor.size = {1, 1, 1};
+    textureDescriptor.sampleCount = 1;
+    textureDescriptor.mipLevelCount = 1;
+    textureDescriptor.usage = wgpu::TextureUsage::RenderAttachment | wgpu::TextureUsage::CopySrc;
+
+    // Wrap ioSurface
+    wgpu::Texture ioSurfaceTexture = WrapIOSurface(&textureDescriptor, ioSurface.get(), true);
+
+    // Uninitialize the teuxture with a render pass.
+    utils::ComboRenderPassDescriptor renderPassDescriptor({ioSurfaceTexture.CreateView()});
+    renderPassDescriptor.cColorAttachments[0].storeOp = wgpu::StoreOp::Discard;
+    wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+    encoder.BeginRenderPass(&renderPassDescriptor).End();
+    wgpu::CommandBuffer commandBuffer = encoder.Finish();
+    queue.Submit(1, &commandBuffer);
+
+    dawn::native::metal::ExternalImageIOSurfaceEndAccessDescriptor endAccessDesc;
+    dawn::native::metal::IOSurfaceEndAccess(ioSurfaceTexture.Get(), &endAccessDesc);
+    EXPECT_FALSE(endAccessDesc.isInitialized);
+}
+
+// Test that an IOSurface may be imported across multiple devices.
+TEST_P(IOSurfaceUsageTests, WriteThenConcurrentReadThenWrite) {
+    DAWN_TEST_UNSUPPORTED_IF(UsesWire());
+
+    ScopedIOSurfaceRef ioSurface = CreateSinglePlaneIOSurface(1, 1, kCVPixelFormatType_32RGBA, 4);
+    uint32_t data = 0x04030201;
+
+    IOSurfaceLock(ioSurface.get(), 0, nullptr);
+    memcpy(IOSurfaceGetBaseAddress(ioSurface.get()), &data, sizeof(data));
+    IOSurfaceUnlock(ioSurface.get(), 0, nullptr);
+
+    // Make additional devices. We will import with the writeDevice, then read it concurrently with
+    // both readDevices.
+    wgpu::Device writeDevice = device;
+    wgpu::Device readDevice1 = CreateDevice();
+    wgpu::Device readDevice2 = CreateDevice();
+
+    wgpu::TextureDescriptor textureDesc;
+    textureDesc.dimension = wgpu::TextureDimension::e2D;
+    textureDesc.format = wgpu::TextureFormat::RGBA8Unorm;
+    textureDesc.size = {1, 1, 1};
+    textureDesc.sampleCount = 1;
+    textureDesc.mipLevelCount = 1;
+    textureDesc.usage = wgpu::TextureUsage::RenderAttachment | wgpu::TextureUsage::CopySrc;
+
+    // Wrap ioSurface
+    dawn::native::metal::ExternalImageDescriptorIOSurface writeExternDesc;
+    writeExternDesc.cTextureDescriptor =
+        reinterpret_cast<const WGPUTextureDescriptor*>(&textureDesc);
+    writeExternDesc.ioSurface = ioSurface.get();
+    writeExternDesc.isInitialized = true;
+
+    wgpu::Texture writeTexture = wgpu::Texture::Acquire(
+        dawn::native::metal::WrapIOSurface(writeDevice.Get(), &writeExternDesc));
+
+    // Clear the texture to green.
+    {
+        utils::ComboRenderPassDescriptor renderPassDescriptor({writeTexture.CreateView()});
+        renderPassDescriptor.cColorAttachments[0].clearValue = {0.0, 1.0, 0.0, 1.0};
+        wgpu::CommandEncoder encoder = writeDevice.CreateCommandEncoder();
+        encoder.BeginRenderPass(&renderPassDescriptor).End();
+        wgpu::CommandBuffer commandBuffer = encoder.Finish();
+        writeDevice.GetQueue().Submit(1, &commandBuffer);
+    }
+
+    // End access.
+    dawn::native::metal::ExternalImageIOSurfaceEndAccessDescriptor endWriteAccessDesc;
+    dawn::native::metal::IOSurfaceEndAccess(writeTexture.Get(), &endWriteAccessDesc);
+    EXPECT_TRUE(endWriteAccessDesc.isInitialized);
+
+    dawn::native::metal::ExternalImageDescriptorIOSurface externDesc;
+    externDesc.cTextureDescriptor = reinterpret_cast<const WGPUTextureDescriptor*>(&textureDesc);
+    externDesc.ioSurface = ioSurface.get();
+    externDesc.isInitialized = true;
+    externDesc.waitEvents.push_back(
+        {endWriteAccessDesc.sharedEvent, endWriteAccessDesc.signaledValue});
+
+    // Wrap on two separate devices to read it.
+    wgpu::Texture readTexture1 =
+        wgpu::Texture::Acquire(dawn::native::metal::WrapIOSurface(readDevice1.Get(), &externDesc));
+    wgpu::Texture readTexture2 =
+        wgpu::Texture::Acquire(dawn::native::metal::WrapIOSurface(readDevice2.Get(), &externDesc));
+
+    // Expect the texture to be green
+    EXPECT_TEXTURE_EQ(readDevice1, utils::RGBA8(0, 255, 0, 255), readTexture1, {0, 0});
+    EXPECT_TEXTURE_EQ(readDevice2, utils::RGBA8(0, 255, 0, 255), readTexture2, {0, 0});
+
+    // End access on both read textures.
+    dawn::native::metal::ExternalImageIOSurfaceEndAccessDescriptor endReadAccessDesc1;
+    dawn::native::metal::IOSurfaceEndAccess(readTexture1.Get(), &endReadAccessDesc1);
+    EXPECT_TRUE(endReadAccessDesc1.isInitialized);
+
+    dawn::native::metal::ExternalImageIOSurfaceEndAccessDescriptor endReadAccessDesc2;
+    dawn::native::metal::IOSurfaceEndAccess(readTexture2.Get(), &endReadAccessDesc2);
+    EXPECT_TRUE(endReadAccessDesc2.isInitialized);
+
+    // Import again for writing. It should not race with the previous reads.
+    writeExternDesc.waitEvents = {endReadAccessDesc1, endReadAccessDesc2};
+    writeExternDesc.isInitialized = true;
+    writeTexture = wgpu::Texture::Acquire(
+        dawn::native::metal::WrapIOSurface(writeDevice.Get(), &writeExternDesc));
+
+    // Clear the texture to blue.
+    {
+        utils::ComboRenderPassDescriptor renderPassDescriptor({writeTexture.CreateView()});
+        renderPassDescriptor.cColorAttachments[0].clearValue = {0.0, 0.0, 1.0, 1.0};
+        wgpu::CommandEncoder encoder = writeDevice.CreateCommandEncoder();
+        encoder.BeginRenderPass(&renderPassDescriptor).End();
+        wgpu::CommandBuffer commandBuffer = encoder.Finish();
+        writeDevice.GetQueue().Submit(1, &commandBuffer);
+    }
+    // Finally, expect the contents to be blue now.
+    EXPECT_TEXTURE_EQ(writeDevice, utils::RGBA8(0, 0, 255, 255), writeTexture, {0, 0});
+    dawn::native::metal::IOSurfaceEndAccess(writeTexture.Get(), &endWriteAccessDesc);
+    EXPECT_TRUE(endWriteAccessDesc.isInitialized);
 }
 
 DAWN_INSTANTIATE_TEST(IOSurfaceValidationTests, MetalBackend());