同步Metal中的网格中的所有线程

我正在尝试为Metal中的n大小的向量编写一个标准或平方长度函数。为此，我计划让每个线程平均每个元素，然后选择一个线程来总和所有元素。同步Metal中的网格中的所有线程

这里是我当前的内核：

#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没有为我实现完全同步。

我是否正确接近这个问题？什么是同步此操作的最佳方式？

来源

2016-12-04 Tyler Hartwig

'simdgroup_barrier'目前仅在iOS上可用。 – warrenm

SIMD组小于线程组，因此同步不起作用。

相反，您需要使用parallel reduction来并行总结这些值。 Here是我找到的一些金属代码。但是，如果您不介意单个线程进行所有求和，您可以使用一个线程运行单独的内核来完成求和。当然，这可能非常缓慢。

来源

2017-06-20 18:48:32 Taylor

同步Metal中的网格中的所有线程

回答

相关问题