オーバーヘッドを実際のスレッドワークよりも高価にすることはできますか？

Question

最近、私はハード最適化の可能性について考えています。つまり、何かを得るためにループを3回反復してハードコードするときに最適化を行うということです。

だから私の考えが一気に現れました。 1024要素のバッファがあるとします。私たちはそれのすべての要素に2を掛けたいと思います。そして、単純なカーネルを作成します。ここでバッファ、outBuffer、サイズ（境界外かどうかを調べるため）と[[thread_position_in_grid]]を渡します。次に単純な乗算を行い、その数を別のバッファに書き出します。

それはそのようなビットになります：

kernel void multiplyBy2(constant float* in [[buffer(0)]], 
          device float* out [[buffer(1)]], 
          constant Uniforms& uniforms [[buffer(2)]], 
          uint gid [[thread_position_in_grid]]) 
{ 

    if (gid >= uniforms.buffer_size) { return; } 
    out[gid] = in[gid] * 2.0; 
} 

、それによって生成され、まだ実際のスレッドの仕事の価値オーバーヘッドが派遣だ場合、私は心配です事はありますか？

それがより効果的であろう、例えば、その

out[gid * 4 + 0] = in[gid + 0] * 2.0; 
    out[gid * 4 + 1] = in[gid + 1] * 2.0; 
    out[gid * 4 + 2] = in[gid + 2] * 2.0; 
    out[gid * 4 + 3] = in[gid + 3] * 2.0; 

ような何かを派遣4倍以下のスレッドが、だからスレッドが少し長く働くことができますか？または、できるだけ細いスレッドを作成する方が良いでしょうか？

Answer 1

はい、これは単なる人為的な例ではなく、実際のシナリオでも同じです。

あなたのような非常に単純なカーネルの場合、ディスパッチのオーバーヘッドは作業を邪魔することがありますが、パフォーマンスにさらに大きな影響を及ぼす要因があります。フェッチされたデータと中間結果の共有。

たとえば、入力テクスチャからピクセルの3x3近傍を読み取り、出力テクスチャに平均値を書き込むカーネルを使用している場合は、フェッチされたテクスチャデータと隣接するピクセル間の部分的な合計を、あなたのカーネル関数内の複数のピクセルとディスパッチするスレッドの総数を減らします。

これはおそらくあなたの好奇心を表します。実際のアプリケーションでは、最適化の前後にすべてのターゲットデバイスでプロファイルする必要があります。