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.
Twoja funkcja kernel jest napisany w sposób całkowicie seryjny, i nie skorzystać z równoległości GPU. Jednak zanim przystąpisz do optymalizacji, jak duży jest twój wektor danych i jak często się zmienia? Jeśli czas przeniesienia danych dominuje nad czasem przetwarzania, użycie GPU może nie być właściwym rozwiązaniem. – warrenm
Tak, jak już podkreślił @warrenm, nie korzystasz z równoległości w GPU. Nie jest tak, ponieważ procesory graficzne działają wydajnie. Musisz wysłać dane do GPU, aby każdy fragment obliczał oddzielny zakres mnożenia. – codetiger
Przykład GPU jest tutaj http://stackoverflow.com/questions/12576976/1d-convolution-without-if-else-statements-non-fft – codetiger