From 270c044100ca5fc5e4b2309997f3dee556386d9e Mon Sep 17 00:00:00 2001 From: Austin Eng Date: Fri, 4 Nov 2022 22:06:20 +0000 Subject: [PATCH] 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 Kokoro: Kokoro Commit-Queue: Austin Eng --- include/dawn/native/MetalBackend.h | 30 ++++ src/dawn/native/metal/CommandBufferMTL.mm | 22 +++ .../native/metal/CommandRecordingContext.h | 5 + src/dawn/native/metal/DeviceMTL.h | 11 +- src/dawn/native/metal/DeviceMTL.mm | 25 +++- src/dawn/native/metal/MetalBackend.mm | 18 ++- src/dawn/native/metal/TextureMTL.h | 13 +- src/dawn/native/metal/TextureMTL.mm | 40 ++++- .../tests/end2end/IOSurfaceWrappingTests.cpp | 137 ++++++++++++++++++ 9 files changed, 287 insertions(+), 14 deletions(-) diff --git a/include/dawn/native/MetalBackend.h b/include/dawn/native/MetalBackend.h index 72a9cb0090..b9ab7070a3 100644 --- a/include/dawn/native/MetalBackend.h +++ b/include/dawn/native/MetalBackend.h @@ -15,6 +15,8 @@ #ifndef INCLUDE_DAWN_NATIVE_METALBACKEND_H_ #define INCLUDE_DAWN_NATIVE_METALBACKEND_H_ +#include + #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`. + // 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 sharedEvent = nil; + static_assert(sizeof(id) == sizeof(void*)); + static_assert(alignof(id) == 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 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 diff --git a/src/dawn/native/metal/CommandBufferMTL.mm b/src/dawn/native/metal/CommandBufferMTL.mm index a5e5fbe52e..2bb83ee1db 100644 --- a/src/dawn/native/metal/CommandBufferMTL.mm +++ b/src/dawn/native/metal/CommandBufferMTL.mm @@ -761,6 +761,10 @@ MaybeError CommandBuffer::FillCommands(CommandRecordingContext* commandContext) case Command::BeginComputePass: { BeginComputePassCmd* cmd = mCommands.NextCommand(); + 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(); + 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, diff --git a/src/dawn/native/metal/CommandRecordingContext.h b/src/dawn/native/metal/CommandRecordingContext.h index ca096bb535..811592c306 100644 --- a/src/dawn/native/metal/CommandRecordingContext.h +++ b/src/dawn/native/metal/CommandRecordingContext.h @@ -22,6 +22,11 @@ namespace dawn::native::metal { +struct MTLSharedEventAndSignalValue { + NSPRef 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 { diff --git a/src/dawn/native/metal/DeviceMTL.h b/src/dawn/native/metal/DeviceMTL.h index fef04bf7a6..f36c22819c 100644 --- a/src/dawn/native/metal/DeviceMTL.h +++ b/src/dawn/native/metal/DeviceMTL.h @@ -18,6 +18,7 @@ #include #include #include +#include #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 CreateTextureWrappingIOSurface(const ExternalImageDescriptor* descriptor, - IOSurfaceRef ioSurface); + void ExportLastSignaledEvent(ExternalImageMTLSharedEventDescriptor* desc); + + Ref CreateTextureWrappingIOSurface( + const ExternalImageDescriptor* descriptor, + IOSurfaceRef ioSurface, + std::vector waitEvents); void WaitForCommandsToBeScheduled(); ResultOrError> CreateStagingBuffer(size_t size) override; @@ -134,6 +140,7 @@ class Device final : public DeviceBase { ResultOrError CheckAndUpdateCompletedSerials() override; NSPRef> mMtlDevice; + NSPRef mMtlSharedEvent = nil; // MTLSharedEvent not available until macOS 10.14+. NSPRef> mCommandQueue; CommandRecordingContext mCommandContext; diff --git a/src/dawn/native/metal/DeviceMTL.mm b/src/dawn/native/metal/DeviceMTL.mm index ee4952ff85..5073e0c6cd 100644 --- a/src/dawn/native/metal/DeviceMTL.mm +++ b/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 sharedEvent = static_cast>(rawEvent); + [*pendingCommands encodeSignalEvent:sharedEvent value:static_cast(pendingSerial)]; + } [*pendingCommands commit]; return mCommandContext.PrepareNextCommandBuffer(*mCommandQueue); } +void Device::ExportLastSignaledEvent(ExternalImageMTLSharedEventDescriptor* desc) { + desc->sharedEvent = *mMtlSharedEvent; + desc->signaledValue = static_cast(GetLastSubmittedCommandSerial()); +} + ResultOrError> Device::CreateStagingBuffer(size_t size) { std::unique_ptr stagingBuffer = std::make_unique(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 Device::CreateTextureWrappingIOSurface(const ExternalImageDescriptor* descriptor, - IOSurfaceRef ioSurface) { +Ref Device::CreateTextureWrappingIOSurface( + const ExternalImageDescriptor* descriptor, + IOSurfaceRef ioSurface, + std::vector waitEvents) { const TextureDescriptor* textureDescriptor = FromAPI(descriptor->cTextureDescriptor); if (ConsumedError(ValidateIsAlive())) { return nullptr; @@ -495,7 +512,9 @@ Ref Device::CreateTextureWrappingIOSurface(const ExternalImageDescripto } Ref result; - if (ConsumedError(Texture::CreateFromIOSurface(this, descriptor, ioSurface), &result)) { + if (ConsumedError( + Texture::CreateFromIOSurface(this, descriptor, ioSurface, std::move(waitEvents)), + &result)) { return nullptr; } return result; diff --git a/src/dawn/native/metal/MetalBackend.mm b/src/dawn/native/metal/MetalBackend.mm index d4140205bb..9bdf138125 100644 --- a/src/dawn/native/metal/MetalBackend.mm +++ b/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 texture = - backendDevice->CreateTextureWrappingIOSurface(cDescriptor, cDescriptor->ioSurface); + std::vector waitEvents; + for (const auto& waitEvent : cDescriptor->waitEvents) { + waitEvents.push_back( + {static_cast>(waitEvent.sharedEvent), waitEvent.signaledValue}); + } + Ref 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(); } diff --git a/src/dawn/native/metal/TextureMTL.h b/src/dawn/native/metal/TextureMTL.h index 3a9c3d8f35..0d5a3c8d0c 100644 --- a/src/dawn/native/metal/TextureMTL.h +++ b/src/dawn/native/metal/TextureMTL.h @@ -17,17 +17,20 @@ #include #import +#include #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> CreateFromIOSurface( Device* device, const ExternalImageDescriptor* descriptor, - IOSurfaceRef ioSurface); + IOSurfaceRef ioSurface, + std::vector waitEvents); static Ref CreateWrapping(Device* device, const TextureDescriptor* descriptor, NSPRef> 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 waitEvents); void InitializeAsWrapping(const TextureDescriptor* descriptor, NSPRef> wrapped); void DestroyImpl() override; @@ -76,6 +84,7 @@ class Texture final : public TextureBase { MTLTextureUsage mMtlUsage; CFRef mIOSurface = nullptr; + std::vector mWaitEvents; }; class TextureView final : public TextureViewBase { diff --git a/src/dawn/native/metal/TextureMTL.mm b/src/dawn/native/metal/TextureMTL.mm index 98d4f3150c..0df13ea02e 100644 --- a/src/dawn/native/metal/TextureMTL.mm +++ b/src/dawn/native/metal/TextureMTL.mm @@ -691,14 +691,17 @@ ResultOrError> Texture::Create(Device* device, const TextureDescrip } // static -ResultOrError> Texture::CreateFromIOSurface(Device* device, - const ExternalImageDescriptor* descriptor, - IOSurfaceRef ioSurface) { +ResultOrError> Texture::CreateFromIOSurface( + Device* device, + const ExternalImageDescriptor* descriptor, + IOSurfaceRef ioSurface, + std::vector waitEvents) { const TextureDescriptor* textureDescriptor = FromAPI(descriptor->cTextureDescriptor); Ref 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 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 sharedEvent = static_cast>(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) {} diff --git a/src/dawn/tests/end2end/IOSurfaceWrappingTests.cpp b/src/dawn/tests/end2end/IOSurfaceWrappingTests.cpp index e9f5b59aa4..6650b74d70 100644 --- a/src/dawn/tests/end2end/IOSurfaceWrappingTests.cpp +++ b/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(&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(&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());