由内核线程递增

Question

假设我想在每次执行内核线程时增加一个属性 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, 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, 0, memory_order_relaxed);

    atomic_fetch_add_explicit(&local_atomic, 1, 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);
    }
}

并按预期工作

Answer 1

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

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, 1, memory_order_relaxed);
    threadgroup_barrier(mem_flags::mem_threadgroup);
    if(lid == 0) {
        atomic_fetch_add_explicit(incremental, local, memory_order_relaxed);
    }
}

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

atomic_fetch_add_explicit 在线程组变量上将使用线程组原子而不是全局原子，后者应该更快。

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

由内核线程递增

Incremental by kernel thread

c++

multithreading

atomic

metal