implements #8119 : Metal backend's wgpu_hal::Device::wait implementation polls instead of waiting

[39ali]

Connections
#8119
#9531

Description
this uses a CondVar and lets the thread sleep instead of polling every 1ms

Testing
ran poll tests

Squash or Rebase?
Rebase

Checklist

• Run cargo fmt.
• Run taplo format.
• Run cargo clippy --tests. If applicable, add:
  • --target wasm32-unknown-unknown
• Run cargo xtask test to run tests.
• If this contains user-facing changes, add a CHANGELOG.md entry.

[39ali]

If timeout is provided, the function will block indefinitely or until , is this a typo ?, should be If timeout is NOT provided

[39ali]

@jimblandy

[inner-daemons]

Looks good overall. 2 questions. Also going to CC @cwfitzgerald because this will almost certainly have to be reassigned

[39ali]

made couple of changes that should address both comments :

• Value inside the mutex (correctness / clarity)
  Previously, completed_value was an AtomicU64 that lived outside the mutex, kinda an anti-pattern, now Mutex<FenceValue> directly guards the value, which is the canonical condvar usage.

• Sync is now per-fence (reduced spurious wakeups)
  Previously sync lived on AdapterShared, shared across every fence submitted to any device from that adapter. When any command buffer completed, it called notify_all() on that shared condvar, waking every thread waiting on any fence from that adapter, even completely unrelated ones. They'd all re-check their predicate and go back to sleep. Now each Fence owns its own condvar, so a completion only wakes threads actually waiting on that fence.

[andyleiserson]

It pains me a little to replace an Atomic with a Mutex, but this is the right way of using the Condvar, and I don't think it's worth speculatively adding a shadow copy in an atomic.

[ruihe774]

Suggestion: add a regression test for #9531.

I'm closing #9532 in favor of this PR. After investigation, the actual failure mode in #9531 is the command buffer terminating in MTLCommandBufferStatusError — the GPU watchdog killing it with kIOGPUCommandBufferCallbackErrorImpactingInteractivity — not just slow successful execution. Your Condvar approach handles this correctly because addCompletedHandler fires on both Completed and Error terminal states, so the wait wakes either way. The MTLSharedEvent approach in #9532 relied on the GPU-encoded encodeSignalEvent_value step, which the GPU never reaches when the CB is killed mid-flight, so it would still deadlock on the original repro.

#9532 included a regression test, which reliably triggers the deadlock on trunk and passes in <1s on this branch:

• this PR + test: pass in 0.79s
• trunk + test: hangs past 90s (TIMEOUT)

I wonder if we can also add this test in this PR. Drop this into tests/tests/wgpu-gpu/poll.rs and add WAIT_INDEFINITELY_LONG_RUNNING to the all_tests list at the top of the file:

/// Regression test for <https://github.com/gfx-rs/wgpu/issues/9531>.
///
/// On Metal, `poll(wait_indefinitely())` deadlocked for command buffers that
/// took more than a few hundred milliseconds because `Device::wait` spin-polled
/// `MTLCommandBuffer.status()` for the `Completed` state. In practice the
/// long-running CB ends up in `MTLCommandBufferStatusError` via the GPU
/// watchdog (`kIOGPUCommandBufferCallbackErrorImpactingInteractivity`), which
/// the spin loop ignored, so the wait never returned.
#[gpu_test]
static WAIT_INDEFINITELY_LONG_RUNNING: GpuTestConfiguration = GpuTestConfiguration::new()
    .parameters(TestParameters::default().test_features_limits())
    .run_async(|ctx| async move {
        // Iteration count tuned so the GPU work is long enough to expose the
        // missed-completion bug in the previous spin-poll implementation
        // (verified to deadlock on trunk on Apple M2 prior to this PR).
        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(wgpu::ShaderModuleDescriptor {
                label: None,
                source: wgpu::ShaderSource::Wgsl(SHADER.into()),
            });

        let buf = ctx.device.create_buffer(&BufferDescriptor {
            label: None,
            size: (N_THREADS as u64) * 4,
            usage: BufferUsages::STORAGE,
            mapped_at_creation: false,
        });

        let bgl = 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(&wgpu::PipelineLayoutDescriptor {
                label: None,
                bind_group_layouts: &[Some(&bgl)],
                immediate_size: 0,
            });

        let pipeline = ctx
            .device
            .create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
                label: None,
                layout: Some(&pipeline_layout),
                module: &module,
                entry_point: Some("main"),
                compilation_options: Default::default(),
                cache: None,
            });

        let bg = ctx.device.create_bind_group(&BindGroupDescriptor {
            label: None,
            layout: &bgl,
            entries: &[BindGroupEntry {
                binding: 0,
                resource: buf.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, &bg, &[]);
            cpass.dispatch_workgroups(N_THREADS / 64, 1, 1);
        }
        ctx.queue.submit(Some(encoder.finish()));

        ctx.async_poll(PollType::wait_indefinitely()).await.unwrap();
    });

[andyleiserson]

The command buffer timeout test could be fragile in CI. One question is whether the paravirtualized GPU device enforces the same "impacting interactivity" timeout that has been observed locally. I'm also not sure if their might be circumstances that cause the effective timeout to vary (which could make the test flaky), or if there's a risk we upset the OS enough to revoke GPU access entirely.

However, I do think a command buffer error test is valuable enough that it's worth at least trying to include in CI. I did look briefly for other ways of exercising the error state for testing, and didn't find anything that seemed better.

[39ali]

i fixed Metal wait on errored command buffers
and added a Metal regression test covering wait_indefinitely on long-running
work that can hit command-buffer error completion.

[andyleiserson]

I added some comments to the test for #9531, revised the changelog entry, and squashed things into a commit for each of the linked bugs so this can be rebase merged.

I also accidentally pushed an old version, then pushed again to undo that, and pushed a third time with the correct changes. https://github.com/gfx-rs/wgpu/compare/a67aa543d0d47b465a3b6edeabf3378fe25dadf0..a408766466a29a71b394a90964da83b822fb6693 shows the actual edits I made. And it looks I will need to push one more time to fix whitespace in the changelog.