2016-12-17 9 views
1
にフロートバッファを充填

問題:メタル

私は一定の値でMTLBufferFloatのSを入力する必要があります - 1729.68921を言います。私はまた、できるだけ早くそれが必要です。

したがって、私はCPU側のバッファを埋めることができません(すなわちUnsafeMutablePointer<Float>MTLBufferから取得し、シリアル方式で割り当てること)。

私のアプローチ

理想的には私はそれがUInt8は、1バイトの長さで、Floatは私がすることはできません、4バイト長であることを考えるとUInt8値(とバッファを埋めることができるだけだが私の知る限り、MTLBlitCommandEncoder.fill()を使用したいです私のFloat定数の任意の値を指定してください)。

これまでのところ、私は唯一の2つのオプションが残って見ることができますが、両方は行き過ぎのように見える:

  1. 一定の値で埋め別のバッファBを作成し、MTLBlitCommandEncoder
  2. 作成を経て、私のバッファにその内容をコピーバッファに

質問

0123を埋めるだろう kernel機能

Floatの最も速い充填方法は、 一定値ですか?

+0

GPUで定数が必要な場合は、その定数をカーネル関数内に作成し、バッファからCPUに送信する代わりに、ローカルで使用するのが最も簡単な方法です。 – Marius

+0

@マリアス私はバッファを埋める必要がある理由を明確にすべきでした:私はものを計算するループを持っています。ループの始めに、バッファをある一定の値に初期化する必要があります。後のループでは、バッファのいくつかの場所*に異なる値を割り当てます。場所はループごとに異なります。 – sarasvati

+1

CPUの使用が遅いのは確実ですか?あなたは 'memset_pattern4()'を知っていますか?それは自分の素朴な実装を超えて最適化される可能性が高いですか? –

答えて

2

各スレッドの複数のバッファ要素に書き込む計算シェーダを使用することは、私の実験で最も速いアプローチでした。これはハードウェアに依存しているため、アプリのデプロイメントが期待されるあらゆるデバイスでテストする必要があります。

私は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倍遅かったです!常に自分の作業負荷でテストしてください。

関連する問題