iOS/Mac OSのメタルコードを高速化するには

Question

長さのある2つのベクトルの間に1D畳み込みを実行するコードをMetalで実装しようとしています。私は正しく

kernel void convolve(const device float *dataVector [[ buffer(0) ]], 
        const device int& dataSize [[ buffer(1) ]], 
        const device float *filterVector [[ buffer(2) ]], 
        const device int& filterSize [[ buffer(3) ]], 
        device float *outVector [[ buffer(4) ]], 
        uint id [[ thread_position_in_grid ]]) { 
    int outputSize = dataSize - filterSize + 1; 
    for (int i=0;i<outputSize;i++) { 
     float sum = 0.0; 
     for (int j=0;j<filterSize;j++) { 
      sum += dataVector[i+j] * filterVector[j]; 
     } 
     outVector[i] = sum; 
    } 
} 

私の問題は、それが（GPUからの/計算+データ転送）を処理するのに約10倍の時間を要している作品次のCPU上でスウィフトよりも金属を使用して同じデータを実装しました。私の質問は、私は内部ループを単一のベクトル演算に置き換えるか、上記のコードを高速化する別の方法はありますか？

Answer 1

この場合、GPUの並列性を利用するための鍵は、外部ループを管理することです。データベクトル全体に対して一度カーネルを呼び出すのではなく、データベクトルの各要素に対してカーネルを呼び出すことになります。カーネル関数は、このように簡略化：この作品を派遣するために

kernel void convolve(const device float *dataVector [[ buffer(0) ]], 
        const constant int &dataSize [[ buffer(1) ]], 
        const constant float *filterVector [[ buffer(2) ]], 
        const constant int &filterSize [[ buffer(3) ]], 
        device float *outVector [[ buffer(4) ]], 
        uint id [[ thread_position_in_grid ]]) 
{ 
    float sum = 0.0; 
    for (int i = 0; i < filterSize; ++i) { 
     sum += dataVector[id + i] * filterVector[i]; 
    } 
    outVector[id] = sum; 
} 

、我々は計算パイプライン状態が推奨するスレッド実行幅に基づいて、スレッドグループサイズを選択します。ここで難しいのは、入力バッファと出力バッファに十分なパディングがあることを確認して、データの実際のサイズを少しオーバーランさせることです。これは、私たちにメモリと計算量を浪費させますが、バッファの終わりにある要素のたたみ込みを計算するために、別のディスパッチを行う複雑さを軽減します。私の実験では

// We should ensure here that the data buffer and output buffer each have a size that is a multiple of 
// the compute pipeline's threadExecutionWidth, by padding the amount we allocate for each of them. 
// After execution, we ignore the extraneous elements in the output buffer beyond the first (dataCount - filterCount + 1). 

let iterationCount = dataCount - filterCount + 1 
let threadsPerThreadgroup = MTLSize(width: min(iterationCount, computePipeline.threadExecutionWidth), height: 1, depth: 1) 
let threadgroups = (iterationCount + threadsPerThreadgroup.width - 1)/threadsPerThreadgroup.width 
let threadgroupsPerGrid = MTLSize(width: threadgroups, height: 1, depth: 1) 

let commandEncoder = commandBuffer.computeCommandEncoder() 
commandEncoder.setComputePipelineState(computePipeline) 
commandEncoder.setBuffer(dataBuffer, offset: 0, at: 0) 
commandEncoder.setBytes(&dataCount, length: MemoryLayout<Int>.stride, at: 1) 
commandEncoder.setBuffer(filterBuffer, offset: 0, at: 2) 
commandEncoder.setBytes(&filterCount, length: MemoryLayout<Int>.stride, at: 3) 
commandEncoder.setBuffer(outBuffer, offset: 0, at: 4) 
commandEncoder.dispatchThreadgroups(threadgroupsPerGrid, threadsPerThreadgroup: threadsPerThreadgroup) 
commandEncoder.endEncoding() 

、この並列化されたアプローチは、問題のシリアルバージョンより400-1000x速く実行されます。私はそれがあなたのCPUの実装とどのように比較されているのか興味があります。