Jak przyspieszyć Metal Code dla iOS/Mac OS

Question

Próbuję zaimplementować kod w metalu, który wykonuje splot 1D między dwoma wektorami o długościach. I zostały wdrożone następujące która działa poprawnie

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

Moim problemem jest to, że trwa około 10 razy dłużej do przetworzenia (transfer danych obliczeniowych + do/z GPU) tych samych danych przy użyciu metalu niż w Swift na CPU. Moje pytanie brzmi: jak zastąpić wewnętrzną pętlę pojedynczą operacją wektorową lub czy istnieje inny sposób na przyspieszenie powyższego kodu?

Answer 1

Kluczem do wykorzystania paralelności GPU w tym przypadku jest umożliwienie zarządzania pętlą zewnętrzną. Zamiast raz wywoływać jądro dla całego wektora danych, wywołajmy go dla każdego elementu w wektorze danych. Funkcja jądra upraszcza to:

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

W celu wysłaniu tej pracy, możemy wybrać rozmiar threadgroup oparciu o szerokości wykonania gwint zalecanej przez państwo rurociągu obliczeniowej. Jedną z najtrudniejszych rzeczy jest upewnienie się, że w buforze wejściowym i wyjściowym jest wystarczająca ilość podkładek, abyśmy mogli nieco przekroczyć rzeczywisty rozmiar danych. To powoduje, że marnujemy niewielką ilość pamięci i obliczeń, ale oszczędzamy nam złożoności wykonywania oddzielnej wysyłki tylko po to, aby obliczyć splot dla elementów na końcu bufora.

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

W moich doświadczeniach, to parallelized podejście prowadzi 400-1000x szybciej niż w wersji seryjnej w pytaniu. Jestem ciekawy, jak to się ma do twojej implementacji procesora.