iOS/Mac OS için Metal Kod Nasıl Hızlandırılır

Question

Metalde kod uygulayarak, iki vektör arasında uzunlukları olan bir 1D dönüşümü gerçekleştiriyorum. Ben

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; 
    } 
} 

Sorunum işlemek için yaklaşık 10 kat daha uzun sürer düzgün çalışıyorsa aşağıdaki uyguladık (hesaplama + veri transferi için/GPU itibaren) bir CPU üzerine Swift daha Metal kullanarak aynı veri. Sorum şu: iç döngüyü tek bir vektör işlemiyle nasıl değiştirebilirim veya yukarıdaki kodu hızlandırmak için başka bir yol var mı?

Answer 1

Bu durumda GPU'nun paralelliklerinden yararlanmanın anahtarı, dış döngüyü sizin için yönetmesini sağlamaktır. Tüm veri vektörü için çekirdeği bir kez çağırmak yerine, veri vektöründeki her eleman için onu çağırırız. Çekirdek fonksiyonu buna kolaylaştırır: Bu çalışmayı sevk etmek

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; 
} 

, biz hesaplama boru hattı devlet tarafından önerilen iplik yürütme genişliğine dayanan bir threadgroup boyutu seçin. Buradaki en zor şey, giriş ve çıkış arabelleklerinde yeterli miktarda dolgu olduğundan emin olmaktır, böylece verilerin gerçek boyutunu biraz aşabiliriz. Bu bize az miktarda bellek ve hesaplama israf etmemize neden olur, ancak bize tamponun sonunda bulunan elemanlar için konvolusyonu hesaplamak için ayrı bir gönderim yapmanın karmaşıklığını korur. Benim deneylerde

// 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() 

, bu parallelized yaklaşım söz konusu seri sürümüne göre daha400-1000x hızlı çalışır. CPU uygulamanızla nasıl kıyaslandığını merak ediyorum.