Fix Metal wait on errored command buffers

and
Add a Metal regression test covering wait_indefinitely on long-running
work that can hit command-buffer error completion.

Author: 39ali

CHANGELOG.md

@@ -236,6 +236,7 @@ By @beholdnec in [#8505](https://github.com/gfx-rs/wgpu/pull/8505).
 #### Metal
 
 - Fix crash on fence creation when running in a MacOS Seatbelt sandbox. By @wumpf in [#9415](https://github.com/gfx-rs/wgpu/pull/9415)
+- Fixed `Device::poll(PollType::wait_indefinitely())` hanging when a Metal command buffer exits with an error. By @39ali in [#9328](https://github.com/gfx-rs/wgpu/pull/9328).
 - Fixed structure field names incorrectly ignoring reserved keywords in the Metal (MSL) backend. By @39ali [#9379](https://github.com/gfx-rs/wgpu/pull/9379).
 - Restore the `Queue::as_raw` method, which was removed without good reason in v29. It now returns `&ProtocolObject<dyn MTLCommandQueue>`. By @andyleiserson in [#9560](https://github.com/gfx-rs/wgpu/pull/9560).
 

tests/tests/wgpu-gpu/poll.rs

@@ -1,18 +1,20 @@
 use std::{num::NonZeroU64, time::Duration};
 
 use wgpu::{
-    BindGroupDescriptor, BindGroupEntry, BindGroupLayoutDescriptor, BindGroupLayoutEntry,
+    Backends, BindGroupDescriptor, BindGroupEntry, BindGroupLayoutDescriptor, BindGroupLayoutEntry,
     BindingResource, BindingType, BufferBindingType, BufferDescriptor, BufferUsages, CommandBuffer,
-    CommandEncoderDescriptor, ComputePassDescriptor, PollType, ShaderStages,
+    CommandEncoderDescriptor, ComputePassDescriptor, ComputePipelineDescriptor,
+    PipelineLayoutDescriptor, PollType, ShaderModuleDescriptor, ShaderSource, ShaderStages,
 };
 
 use wgpu_test::{
-    gpu_test, GpuTestConfiguration, GpuTestInitializer, TestParameters, TestingContext,
+    gpu_test, FailureCase, GpuTestConfiguration, GpuTestInitializer, TestParameters, TestingContext,
 };
 
 pub fn all_tests(vec: &mut Vec<GpuTestInitializer>) {
     vec.extend([
         WAIT,
+        WAIT_INDEFINITELY_LONG_RUNNING,
         WAIT_WITH_TIMEOUT,
         WAIT_WITH_TIMEOUT_MAX,
         DOUBLE_WAIT,
@@ -85,6 +87,95 @@ static WAIT: GpuTestConfiguration = GpuTestConfiguration::new()
         .unwrap();
     });
 
+/// Regression test for <https://github.com/gfx-rs/wgpu/issues/9531>.
+#[gpu_test]
+static WAIT_INDEFINITELY_LONG_RUNNING: GpuTestConfiguration = GpuTestConfiguration::new()
+    .parameters(
+        TestParameters::default()
+            .test_features_limits()
+            .skip(FailureCase::backend(!Backends::METAL)),
+    )
+    .run_async(|ctx| async move {
+        const SHADER: &str = r#"
+@group(0) @binding(0) var<storage, read_write> buf: array<u32>;
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+    var x: u32 = gid.x ^ 0xDEADBEEFu;
+    for (var i: u32 = 0u; i < 5000000u; i++) {
+        x ^= x << 13u;
+        x ^= x >> 17u;
+        x ^= x << 5u;
+    }
+    buf[gid.x] = x;
+}
+"#;
+
+        const N_THREADS: u32 = 1024 * 64;
+
+        let module = ctx.device.create_shader_module(ShaderModuleDescriptor {
+            label: None,
+            source: ShaderSource::Wgsl(SHADER.into()),
+        });
+        let buffer = ctx.device.create_buffer(&BufferDescriptor {
+            label: None,
+            size: (N_THREADS as u64) * 4,
+            usage: BufferUsages::STORAGE,
+            mapped_at_creation: false,
+        });
+        let bind_group_layout = ctx
+            .device
+            .create_bind_group_layout(&BindGroupLayoutDescriptor {
+                label: None,
+                entries: &[BindGroupLayoutEntry {
+                    binding: 0,
+                    visibility: ShaderStages::COMPUTE,
+                    ty: BindingType::Buffer {
+                        ty: BufferBindingType::Storage { read_only: false },
+                        has_dynamic_offset: false,
+                        min_binding_size: None,
+                    },
+                    count: None,
+                }],
+            });
+        let pipeline_layout = ctx
+            .device
+            .create_pipeline_layout(&PipelineLayoutDescriptor {
+                label: None,
+                bind_group_layouts: &[Some(&bind_group_layout)],
+                immediate_size: 0,
+            });
+        let pipeline = ctx
+            .device
+            .create_compute_pipeline(&ComputePipelineDescriptor {
+                label: None,
+                layout: Some(&pipeline_layout),
+                module: &module,
+                entry_point: Some("main"),
+                compilation_options: Default::default(),
+                cache: None,
+            });
+        let bind_group = ctx.device.create_bind_group(&BindGroupDescriptor {
+            label: None,
+            layout: &bind_group_layout,
+            entries: &[BindGroupEntry {
+                binding: 0,
+                resource: buffer.as_entire_binding(),
+            }],
+        });
+        let mut encoder = ctx
+            .device
+            .create_command_encoder(&CommandEncoderDescriptor::default());
+        {
+            let mut cpass = encoder.begin_compute_pass(&ComputePassDescriptor::default());
+            cpass.set_pipeline(&pipeline);
+            cpass.set_bind_group(0, &bind_group, &[]);
+            cpass.dispatch_workgroups(N_THREADS / 64, 1, 1);
+        }
+        ctx.queue.submit(Some(encoder.finish()));
+        ctx.async_poll(PollType::wait_indefinitely()).await.unwrap();
+    });
+
 #[gpu_test]
 static WAIT_WITH_TIMEOUT: GpuTestConfiguration = GpuTestConfiguration::new()
     .parameters(TestParameters::default().enable_noop())

wgpu-hal/src/metal/device.rs

@@ -1922,19 +1922,13 @@ impl crate::Device for super::Device {
             }
         }
 
-        if let Some(deadline) =
-            timeout.and_then(|timeout| std::time::Instant::now().checked_add(timeout))
-        {
-            while *lock < wait_value {
-                let result = condvar.wait_until(&mut lock, deadline);
-                if result.timed_out() {
-                    return Ok(*lock >= wait_value);
-                }
+        if let Some(timeout) = timeout {
+            let result = condvar.wait_while_for(&mut lock, |value| *value < wait_value, timeout);
+            if result.timed_out() {
+                return Ok(*lock >= wait_value);
             }
         } else {
-            while *lock < wait_value {
-                condvar.wait(&mut lock);
-            }
+            condvar.wait_while(&mut lock, |value| *value < wait_value);
         }
 
         Ok(true)

wgpu-hal/src/metal/mod.rs

@@ -1050,8 +1050,11 @@ impl Fence {
         let mut max_value = *self.sync.0.lock();
         let pending_command_buffers = self.pending_command_buffers.read();
         for &(value, ref cmd_buf) in pending_command_buffers.iter() {
-            if cmd_buf.status() == MTLCommandBufferStatus::Completed {
-                max_value = value;
+            match cmd_buf.status() {
+                MTLCommandBufferStatus::Completed | MTLCommandBufferStatus::Error => {
+                    max_value = value;
+                }
+                _ => {}
             }
         }
         max_value