Быстрый параллельный расчет суммы массива на iOS

Question

На основании @Kametrixom answer, я сделал несколько тестовых приложений для параллельного вычисления суммы в массиве.

Моя тестовая программа выглядит следующим образом:

import UIKit 
import Metal 

class ViewController: UIViewController { 

// Data type, has to be the same as in the shader 
typealias DataType = CInt 

override func viewDidLoad() { 
    super.viewDidLoad() 

    let data = (0..<10000000).map{ _ in DataType(200) } // Our data, randomly generated 


    var start, end : UInt64 


    var result:DataType = 0 
    start = mach_absolute_time() 
    data.withUnsafeBufferPointer { buffer in 
     for elem in buffer { 
      result += elem 
     } 
    } 
    end = mach_absolute_time() 

    print("CPU result: \(result), time: \(Double(end - start)/Double(NSEC_PER_SEC))") 

    result = 0 


    start = mach_absolute_time() 
    result = sumParallel4(data) 
    end = mach_absolute_time() 

    print("Metal result: \(result), time: \(Double(end - start)/Double(NSEC_PER_SEC))") 


    result = 0 

    start = mach_absolute_time() 
    result = sumParralel(data) 
    end = mach_absolute_time() 

    print("Metal result: \(result), time: \(Double(end - start)/Double(NSEC_PER_SEC))") 

    result = 0 

    start = mach_absolute_time() 
    result = sumParallel3(data) 
    end = mach_absolute_time() 

    print("Metal result: \(result), time: \(Double(end - start)/Double(NSEC_PER_SEC))") 





} 

func sumParralel(data : Array<DataType>) -> DataType { 

    let count = data.count 
    let elementsPerSum: Int = Int(sqrt(Double(count))) 

    let device = MTLCreateSystemDefaultDevice()! 
    let parsum = device.newDefaultLibrary()!.newFunctionWithName("parsum")! 
    let pipeline = try! device.newComputePipelineStateWithFunction(parsum) 


    var dataCount = CUnsignedInt(count) 
    var elementsPerSumC = CUnsignedInt(elementsPerSum) 
    let resultsCount = (count + elementsPerSum - 1)/elementsPerSum // Number of individual results = count/elementsPerSum (rounded up) 


    let dataBuffer = device.newBufferWithBytes(data, length: strideof(DataType) * count, options: []) // Our data in a buffer (copied) 
    let resultsBuffer = device.newBufferWithLength(strideof(DataType) * resultsCount, options: []) // A buffer for individual results (zero initialized) 
    let results = UnsafeBufferPointer<DataType>(start: UnsafePointer(resultsBuffer.contents()), count: resultsCount) // Our results in convenient form to compute the actual result later 

    let queue = device.newCommandQueue() 
    let cmds = queue.commandBuffer() 
    let encoder = cmds.computeCommandEncoder() 

    encoder.setComputePipelineState(pipeline) 

    encoder.setBuffer(dataBuffer, offset: 0, atIndex: 0) 
    encoder.setBytes(&dataCount, length: sizeofValue(dataCount), atIndex: 1) 
    encoder.setBuffer(resultsBuffer, offset: 0, atIndex: 2) 
    encoder.setBytes(&elementsPerSumC, length: sizeofValue(elementsPerSumC), atIndex: 3) 

    // We have to calculate the sum `resultCount` times => amount of threadgroups is `resultsCount`/`threadExecutionWidth` (rounded up) because each threadgroup will process `threadExecutionWidth` threads 
    let threadgroupsPerGrid = MTLSize(width: (resultsCount + pipeline.threadExecutionWidth - 1)/pipeline.threadExecutionWidth, height: 1, depth: 1) 

    // Here we set that each threadgroup should process `threadExecutionWidth` threads, the only important thing for performance is that this number is a multiple of `threadExecutionWidth` (here 1 times) 
    let threadsPerThreadgroup = MTLSize(width: pipeline.threadExecutionWidth, height: 1, depth: 1) 

    encoder.dispatchThreadgroups(threadgroupsPerGrid, threadsPerThreadgroup: threadsPerThreadgroup) 
    encoder.endEncoding() 


    var result : DataType = 0 


    cmds.commit() 
    cmds.waitUntilCompleted() 
    for elem in results { 
     result += elem 
    } 


    return result 
} 



func sumParralel1(data : Array<DataType>) -> UnsafeBufferPointer<DataType> { 

    let count = data.count 
    let elementsPerSum: Int = Int(sqrt(Double(count))) 

    let device = MTLCreateSystemDefaultDevice()! 
    let parsum = device.newDefaultLibrary()!.newFunctionWithName("parsum")! 
    let pipeline = try! device.newComputePipelineStateWithFunction(parsum) 


    var dataCount = CUnsignedInt(count) 
    var elementsPerSumC = CUnsignedInt(elementsPerSum) 
    let resultsCount = (count + elementsPerSum - 1)/elementsPerSum // Number of individual results = count/elementsPerSum (rounded up) 

    let dataBuffer = device.newBufferWithBytes(data, length: strideof(DataType) * count, options: []) // Our data in a buffer (copied) 
    let resultsBuffer = device.newBufferWithLength(strideof(DataType) * resultsCount, options: []) // A buffer for individual results (zero initialized) 
    let results = UnsafeBufferPointer<DataType>(start: UnsafePointer(resultsBuffer.contents()), count: resultsCount) // Our results in convenient form to compute the actual result later 

    let queue = device.newCommandQueue() 
    let cmds = queue.commandBuffer() 
    let encoder = cmds.computeCommandEncoder() 

    encoder.setComputePipelineState(pipeline) 

    encoder.setBuffer(dataBuffer, offset: 0, atIndex: 0) 
    encoder.setBytes(&dataCount, length: sizeofValue(dataCount), atIndex: 1) 
    encoder.setBuffer(resultsBuffer, offset: 0, atIndex: 2) 
    encoder.setBytes(&elementsPerSumC, length: sizeofValue(elementsPerSumC), atIndex: 3) 

    // We have to calculate the sum `resultCount` times => amount of threadgroups is `resultsCount`/`threadExecutionWidth` (rounded up) because each threadgroup will process `threadExecutionWidth` threads 
    let threadgroupsPerGrid = MTLSize(width: (resultsCount + pipeline.threadExecutionWidth - 1)/pipeline.threadExecutionWidth, height: 1, depth: 1) 

    // Here we set that each threadgroup should process `threadExecutionWidth` threads, the only important thing for performance is that this number is a multiple of `threadExecutionWidth` (here 1 times) 
    let threadsPerThreadgroup = MTLSize(width: pipeline.threadExecutionWidth, height: 1, depth: 1) 

    encoder.dispatchThreadgroups(threadgroupsPerGrid, threadsPerThreadgroup: threadsPerThreadgroup) 
    encoder.endEncoding() 


    cmds.commit() 
    cmds.waitUntilCompleted() 



    return results 
} 

func sumParallel3(data : Array<DataType>) -> DataType { 

    var results = sumParralel1(data) 

    repeat { 
     results = sumParralel1(Array(results)) 
    } while results.count >= 100 

    var result : DataType = 0 

    for elem in results { 
     result += elem 
    } 


    return result 
} 

func sumParallel4(data : Array<DataType>) -> DataType { 

    let queue = NSOperationQueue() 
    queue.maxConcurrentOperationCount = 4 

    var a0 : DataType = 0 
    var a1 : DataType = 0 
    var a2 : DataType = 0 
    var a3 : DataType = 0 

    let op0 = NSBlockOperation(block : { 

     for i in 0..<(data.count/4) { 
      a0 = a0 + data[i] 
     } 

    }) 

    let op1 = NSBlockOperation(block : { 
     for i in (data.count/4)..<(data.count/2) { 
      a1 = a1 + data[i] 
     } 
    }) 

    let op2 = NSBlockOperation(block : { 
     for i in (data.count/2)..<(3 * data.count/4) { 
      a2 = a2 + data[i] 
     } 
    }) 

    let op3 = NSBlockOperation(block : { 
     for i in (3 * data.count/4)..<(data.count) { 
      a3 = a3 + data[i] 
     } 
    }) 



    queue.addOperation(op0) 
    queue.addOperation(op1) 
    queue.addOperation(op2) 
    queue.addOperation(op3) 

    queue.suspended = false 
    queue.waitUntilAllOperationsAreFinished() 

    let aaa: DataType = a0 + a1 + a2 + a3 

    return aaa 
} 
} 

И у меня есть шейдер, который выглядит следующим образом:

kernel void parsum(const device DataType* data [[ buffer(0) ]], 
       const device uint& dataLength [[ buffer(1) ]], 
       device DataType* sums [[ buffer(2) ]], 
       const device uint& elementsPerSum [[ buffer(3) ]], 

       const uint tgPos [[ threadgroup_position_in_grid ]], 
       const uint tPerTg [[ threads_per_threadgroup ]], 
       const uint tPos [[ thread_position_in_threadgroup ]]) { 

    uint resultIndex = tgPos * tPerTg + tPos; // This is the index of the individual result, this var is unique to this thread 
    uint dataIndex = resultIndex * elementsPerSum; // Where the summation should begin 
    uint endIndex = dataIndex + elementsPerSum < dataLength ? dataIndex + elementsPerSum : dataLength; // The index where summation should end 

    for (; dataIndex < endIndex; dataIndex++) 
     sums[resultIndex] += data[dataIndex]; 
} 

На моей неожиданной функции sumParallel4 является самым быстрым, что я думал, что это не должен» t be. Я заметил, что когда я вызываю функции sumParralel и sumParallel3, первая функция всегда медленнее, даже если я изменяю порядок функции. (Так что если я сначала вызову sumParralel, это будет медленнее, если я вызову sumParallel3, это будет медленнее.).

Почему это? Почему sumParallel3 не намного быстрее, чем sumParallel? Почему sumParallel4 самый быстрый, хотя он рассчитан на CPU?

---

Как я могу обновить функцию своего GPU с помощью posix_memalign? Я знаю, что он должен работать быстрее, потому что он имел бы общую память между GPU и CPU, но я не знаю, что массив ведьм должен быть выделен таким образом (данные или результат) и как я могу распределять данные с помощью posix_memalign, если данные передаются параметром ?

Answer 1

При выполнении этих тестов на iPhone 6 я видел, что версия Metal работает между 3x медленнее и 2x быстрее, чем наивное суммирование процессора. С изменениями, которые я описываю ниже, он был последовательно быстрее.

Я обнаружил, что большая часть затрат на запуск версии Metal может быть отнесена не только к распределению буферов, хотя это и было значительным, но и к первому созданию устройства и вычислению состояния конвейера. Это действия, которые вы обычно выполняете один раз при инициализации приложения, поэтому нецелесообразно включать их в синхронизацию.

Следует также отметить, что если вы используете эти тесты с помощью Xcode с включенным уровнем проверки металла и захватом кадров графического процессора, это имеет значительную временную стоимость и искажает результаты в пользу процессора.

С этими предостережениями, вот как вы можете использовать posix_memalign для выделения памяти, которая может использоваться для поддержки MTLBuffer. Хитрость заключается в том, чтобы гарантировать, что запрошенная вами память фактически выровнена по странице (то есть ее адрес кратен getpagesize()), что может привести к округлению объема памяти, превышающей объем фактической необходимости хранения ваших данных:

let dataCount = 1_000_000 
let dataSize = dataCount * strideof(DataType) 
let pageSize = Int(getpagesize()) 
let pageCount = (dataSize + (pageSize - 1))/pageSize 
var dataPointer: UnsafeMutablePointer<Void> = nil 
posix_memalign(&dataPointer, pageSize, pageCount * pageSize) 
let data = UnsafeMutableBufferPointer(start: UnsafeMutablePointer<DataType>(dataPointer), 
             count: (pageCount * pageSize)/strideof(DataType)) 

for i in 0..<dataCount { 
    data[i] = 200 
} 

Это требует создания data в UnsafeMutableBufferPointer<DataType>, а не [DataType], поскольку Свифта Array выделяет свой бэк-магазин. Вам также нужно будет передать количество элементов данных для работы, так как измененного указателя буфера округлены, чтобы выравнивать страницу с буфером.

Чтобы создать MTLBuffer с этими данными, используйте API newBufferWithBytesNoCopy(_:length:options:deallocator:).Крайне важно, чтобы еще раз указанная длина была кратной размеру страницы; в противном случае этот метод возвращает nil:

let roundedUpDataSize = strideof(DataType) * data.count 
let dataBuffer = device.newBufferWithBytesNoCopy(data.baseAddress, length: roundedUpDataSize, options: [], deallocator: nil) 

Здесь мы не предоставляем deallocator, но вы должны освободить память, когда вы закончите его использования, путем пропускания baseAddress указателя буфера для free().