各スレッドの複数のバッファ要素に書き込む計算シェーダを使用することは、私の実験で最も速いアプローチでした。これはハードウェアに依存しているため、アプリのデプロイメントが期待されるあらゆるデバイスでテストする必要があります。
私は2つのコンピュートシェーダ書いた:あなたの場合は
kernel void fill_16_unchecked(device float *buffer [[buffer(0)]],
constant float &value [[buffer(1)]],
uint index [[thread_position_in_grid]])
{
for (int i = 0; i < 16; ++i) {
buffer[index * 16 + i] = value;
}
}
kernel void single_fill_checked(device float *buffer [[buffer(0)]],
constant float &value [[buffer(1)]],
constant uint &buffer_length [[buffer(2)]],
uint index [[thread_position_in_grid]])
{
if (index < buffer_length) {
buffer[index] = value;
}
}
:配列境界に対してチェックせずに16個の連続する配列要素を満たし1、およびバッファの長さに対してチェックした後、単一の配列要素を設定し、1をあなたのバッファ数はスレッド実行幅にループ内で設定した要素の数を掛けたものの倍数になることを知っていれば、最初の関数を使うことができます。 2番目の関数は、そうでなければバッファをオーバーランさせるグリッドをディスパッチする場合の代替手段です。
あなたはこれらの関数から構築された2つのパイプラインを持っていたら、次のように、あなたは計算コマンドのペアで仕事を派遣することができます
NSInteger executionWidth = [unchecked16Pipeline threadExecutionWidth];
id<MTLComputeCommandEncoder> computeEncoder = [commandBuffer computeCommandEncoder];
[computeEncoder setBuffer:buffer offset:0 atIndex:0];
[computeEncoder setBytes:&value length:sizeof(float) atIndex:1];
if (bufferCount/(executionWidth * 16) != 0) {
[computeEncoder setComputePipelineState:unchecked16Pipeline];
[computeEncoder dispatchThreadgroups:MTLSizeMake(bufferCount/(executionWidth * 16), 1, 1)
threadsPerThreadgroup:MTLSizeMake(executionWidth, 1, 1)];
}
if (bufferCount % (executionWidth * 16) != 0) {
int remainder = bufferCount % (executionWidth * 16);
[computeEncoder setComputePipelineState:checkedSinglePipeline];
[computeEncoder setBytes:&bufferCount length:sizeof(bufferCount) atIndex:2];
[computeEncoder dispatchThreadgroups:MTLSizeMake((remainder/executionWidth) + 1, 1, 1)
threadsPerThreadgroup:MTLSizeMake(executionWidth, 1, 1)];
}
[computeEncoder endEncoding];
このように仕事をしていることは、必ずしもより速くなりません。なおスレッドごとに1つの要素だけを書き込む単純なアプローチです。私のテストでは、A8の方が40%速く、A10とほぼ同等で、A9の方が2-3倍遅かったです!常に自分の作業負荷でテストしてください。
GPUで定数が必要な場合は、その定数をカーネル関数内に作成し、バッファからCPUに送信する代わりに、ローカルで使用するのが最も簡単な方法です。 – Marius
@マリアス私はバッファを埋める必要がある理由を明確にすべきでした:私はものを計算するループを持っています。ループの始めに、バッファをある一定の値に初期化する必要があります。後のループでは、バッファのいくつかの場所*に異なる値を割り当てます。場所はループごとに異なります。 – sarasvati
CPUの使用が遅いのは確実ですか?あなたは 'memset_pattern4()'を知っていますか?それは自分の素朴な実装を超えて最適化される可能性が高いですか? –