メタルのグリッド内のすべてのスレッドを同期する

Question

メタルのnサイズのベクトルにノルムまたは2乗の長さ関数を書き込もうとしています。これを行うために、私はすべてのスレッドが各要素を正方形にし、すべての要素を合計するために1つのスレッドを選択することを計画しました。

#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は私が見つけたいくつかの金属コードです。

すべての集計を行っているスレッドが1つでも気にしない場合は、合計を行うスレッドが1つのみの別のカーネルを実行できます。もちろん、これは非常に遅くなる可能性があります。