如何在Metal计算着色器中基于CAS实现正确的加锁解锁序列

你的代码有两个核心问题，导致其他线程永远拿不到锁：

1. 内存序使用完全错误
   你所有原子操作都用了memory_order_relaxed，这个内存序仅保证单个原子操作本身的原子性，不提供任何跨线程的内存可见性保证。解锁时写入的false值不会被及时同步给其他线程，等待线程读到的永远是GPU缓存中旧的true值，自然永远无法跳出循环。
2. 忙等逻辑不符合GPU执行模型
   Metal作为GPU计算API，是以SIMD组（通常32或64个线程为一组）为单位锁步调度执行的。你写的空转忙等循环会占满SIMD组的执行资源，持有锁的线程如果和等待线程在同一个SIMD组，根本拿不到执行周期去完成临界区操作、执行解锁逻辑，直接触发死锁。

另外你使用的atomic_compare_exchange_weak_explicit允许伪失败（即值匹配时也可能返回失败），虽然你每次循环重置expected的写法不会导致逻辑错误，但会增加不必要的重试开销，锁场景下不推荐使用。

按照以下规则修改即可实现正常工作的CAS自旋锁：

• 替换原子操作内存序：加锁CAS成功时使用memory_order_acquire语义，保证拿到锁后可以看到之前持有锁线程写入的所有数据；解锁store操作使用memory_order_release语义，保证临界区的所有写入操作在解锁前完成，对后续拿到锁的线程可见。
• 替换空转忙等逻辑：不要写无意义的空循环，使用Metal提供的原子等待接口主动让出执行资源，同时强制刷新内存读取锁变量的最新值，避免锁步调度导致的死锁和总线争抢。
• 锁场景优先使用strong版本的CAS操作，减少weak版本伪失败带来的额外开销。

修正后的可运行代码如下：

kernel void testFunction(
    device float *depthBuffer [[buffer(4)]],
    device atomic_bool *depthFlag [[buffer(5)]],
    uint index [[thread_position_in_grid]]
) {
    // 加锁
    bool expected = false;
    while (!atomic_compare_exchange_strong_explicit(
        &depthFlag[1],
        &expected,
        true,
        memory_order_acquire,
        memory_order_relaxed
    )) {
        expected = false;
        // 等待锁变量变化，主动让出执行权，避免空转死锁
        atomic_wait_explicit(&depthFlag[1], true, memory_order_relaxed);
    }

    // 此处编写临界区业务逻辑

    // 解锁
    atomic_store_explicit(&depthFlag[1], false, memory_order_release);
    // 唤醒一个等待锁的线程
    atomic_notify_one(&depthFlag[1]);
}

• 如果你的运行环境不支持atomic_wait_explicit/atomic_notify_one接口（Metal 2.2以下版本），忙等阶段需要每重试若干次主动触发设备内存屏障刷新缓存，同时避免所有等待线程持续争抢原子变量导致的活锁。
• 如果仅需要线程组内的同步，不要使用device地址空间的全局锁，优先使用threadgroup地址空间的原子变量配合threadgroup_barrier，性能会高一个数量级以上。
• GPU上全局锁的开销极高，仅适合保护极小占比的临界区操作，不要把大量逻辑放在锁内，否则会完全抵消GPU的并行计算优势。

内容的提问来源于stack exchange，提问作者Magodbaru