c++ - 按内核线程增量

Question

假设我想在每次执行内核线程时增加一个属性var增量: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)。这使得对增量的所有访问都是原子的。这里的 Memory_order_relaxed 意味着放宽了对操作顺序的保证，因此仅当您只是添加或只是从值中减去时，这才有效。 memory_order_relaxed 是 MSL 中唯一支持的 memory_order。您可以在Metal Shading Language Specification中阅读更多相关内容。 ，第 6.13 节。

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

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);
    }
}

这基本上意味着:线程组中的每个线程都应该自动向 local 添加 1，等待每个线程完成 (threadgroup_barrier)，然后恰好有一个线程自动添加总数本地到增量。

线程组变量上的

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

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