由内核线程递增

如何解决由内核线程递增

假设我想在每次执行内核线程时增加一个属性 var incremental:Int32：

//SWIFT
var incremental:Int32 = 0
var incrementalBuffer:MTLBuffer!
var incrementalPointer: UnsafeMutablePointer<Int32>!

init(metalView: MTKView) {
    ...
    incrementalBuffer = Renderer.device.makeBuffer(bytes: &incremental,length: MemoryLayout<Int32>.stride)
    incrementalPointer = incrementalBuffer.contents().bindMemory(to: Int32.self,capacity: 1)
}
func draw(in view: MTKView) {
    ...
    computeCommandEncoder.setComputePipelineState(computePipelineState)
    let width = computePipelineState.threadExecutionWidth
    let threadsPerGroup = MTLSizeMake(width,1,1)
    let threadsPerGrid = MTLSizeMake(10,1)
    computeCommandEncoder.setBuffer(incrementalBuffer,offset: 0,index: 0)
    computeCommandEncoder.dispatchThreads(threadsPerGrid,threadsPerThreadgroup: threadsPerGroup)
    computeCommandEncoder.endEncoding()
    commandBufferCompute.commit()
    commandBufferCompute.waitUntilCompleted()
    
    print(incrementalPointer.pointee)
}

//METAL
kernel void compute_shader (device int& incremental [[buffer(0)]]){
    incremental++;
}

所以我期待输出：

10
20
30
...

但我明白了：

1
2
3
...

编辑： 在根据@JustSomeGuy、raywenderlich 的 Caroline 和一位 Apple 工程师的回答做了一些工作后，我得到了：

[[kernel]] void compute_shader (device atomic_int& incremental [[buffer(0)]],ushort lid [[thread_position_in_threadgroup]] ){

    threadgroup atomic_int local_atomic;
    if (lid==0) atomic_store_explicit(&local_atomic,memory_order_relaxed);

    atomic_fetch_add_explicit(&local_atomic,memory_order_relaxed);

    threadgroup_barrier(mem_flags::mem_threadgroup);

    if(lid == 0) {
        int local_non_atomic = atomic_load_explicit(&local_atomic,memory_order_relaxed);
        atomic_fetch_add_explicit(&incremental,local_non_atomic,memory_order_relaxed);
    }
}

并按预期工作

解决方法

您看到此问题的原因是 ++ 不是原子的。它基本上归结为这样的代码

auto temp = incremental;
incremental = temp + 1;
temp;

这意味着因为线程是“并行”执行的（这不是真的，因为许多线程形成了一个以步进锁定方式执行的 SIMD 组，但这在这里并不重要）。

由于访问不是原子的，结果基本上是未定义的，因为无法判断哪个线程观察到哪个值。

快速解决方法是使用 atomic_fetch_add_explicit(incremental,1,memory_order_relaxed)。这使得对 incremental 的所有访问都是原子的。 memory_order_relaxed 在这里意味着对操作顺序的保证是放宽的，因此只有当您只是从值中添加或减去时，这才有效。 memory_order_relaxed 是 MSL 中唯一支持的 memory_order。您可以在 Metal Shading Language Specification 第 6.13 节中阅读更多相关内容。

但是这个快速修复非常糟糕，因为它会很慢，因为对 incremental 的访问必须在所有线程之间同步。另一种方法是使用一种通用模式，其中线程组中的所有线程更新 threadgroup 内存中的值，然后一个或多个线程自动更新 device 内存。所以内核看起来像

kernel void compute_shader (device int& incremental [[buffer(0)]],threadgroup int& local [[threadgroup(0)]],ushort lid [[thread_position_in_threadgroup]] ){
    atomic_fetch_add_explicit(local,memory_order_relaxed);
    threadgroup_barrier(mem_flags::mem_threadgroup);
    if(lid == 0) {
        atomic_fetch_add_explicit(incremental,local,memory_order_relaxed);
    }
}

这基本上意味着：线程组中的每个线程都应该原子地将 1 添加到 local，等到每个线程完成（threadgroup_barrier），然后正好有一个线程原子地将总 local 添加到incremental。

线程组变量上的

atomic_fetch_add_explicit 将使用线程组原子而不是应该更快的全局原子。

您可以阅读我上面链接的规范以了解更多信息，那里的示例中提到了这些模式。