multithreading - 在Metal中同步网格中的所有线程

Question

我正在尝试为Metal中的n尺寸向量编写一个范数或平方长度函数。为此，我计划让每个线程平方每个元素，然后选择一个线程来求和所有元素。

这是我当前的内核:

#include <metal_stdlib>
#include <metal_compute>
using namespace metal;

kernel void length_squared(const device float *x [[ buffer(0) ]],
                           device float *s [[ buffer(1) ]],
                           device float *out [[ buffer(2) ]],
                           uint gid [[ thread_position_in_grid ]],
                           uint numElements [[ threads_per_grid ]])
{
    s[gid] = x[gid];// * x[gid];
    simdgroup_barrier(mem_flags::mem_none);
    if(gid == 0){
        for(uint i = 0; i < numElements; i++){
            *out += s[i];
        }
    }
}

不幸的是，对于“使用未声明的标识符 simdgroup_barrier”，此代码无法编译。该方法记录在 Metal Shading Language Specification中。

有人遇到过这个吗？或知道如何同步网格中的所有线程？ threadgroup_barrier对我而言未实现完全同步。

我是否错误地解决了这个问题？同步此操作的最佳方法是什么？

Answer 1

SIMD组小于线程组，因此同步将无法进行。

相反，您将需要使用parallel reduction来并行汇总这些值。 Here是我发现的一些 Metal 代码。

但是，如果您不介意由一个线程来进行所有求和，则可以只运行一个线程来运行一个单独的内核来进行求和。当然，这可能非常慢。