Swift Metall parallele Summenberechnung von Array auf iOS

Question

Basierend auf @Kametrixom answer, habe ich einige Test-Anwendung für die parallele Berechnung der Summe in einem Array gemacht.

Meine Testanwendung sieht wie folgt aus:

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

Und ich habe einen Shader, der wie folgt aussieht:

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

Auf meiner Überraschung Funktion sumParallel4 ist das schnellste, was ich dachte, es shouldn‘ t sein. Ich habe festgestellt, dass beim ersten Aufruf der Funktionen sumParralel und sumParallel3 die erste Funktion immer langsamer ist, auch wenn ich die Reihenfolge der Funktionen ändere. (Also, wenn ich sumParralel zuerst anrufe, ist das langsamer, wenn ich sumParallel3 anrufe, ist das langsamer.).

Warum ist das? Warum ist sumParallel3 nicht viel schneller als sumParallel? Warum ist sumParallel4 am schnellsten, obwohl es auf der CPU berechnet wird?

---

Wie kann ich meine GPU Funktion mit posix_memalign aktualisieren? Ich weiß, dass es schneller arbeiten sollte, weil es Speicher zwischen GPU und CPU geteilt hätte, aber ich weiß nicht, dass Hex-Array auf diese Weise zugeordnet werden sollte (Daten oder Ergebnis) und wie kann ich Daten mit posix_memalign zuweisen, wenn Daten in Funktion übergeben werden ?

Answer 1

Beim Ausführen dieser Tests auf einem iPhone 6 sah ich die Metal-Version zwischen 3x langsamer und 2x schneller als die naive CPU-Summierung. Mit den Modifikationen, die ich unten beschreibe, war es immer schneller.

Ich fand heraus, dass ein Großteil der Kosten bei der Ausführung der Metal-Version nicht nur auf die Zuweisung der Puffer zurückzuführen war, obwohl dies signifikant war, sondern auch auf die erstmalige Erstellung des Geräte- und Rechner-Pipeline-Zustands. Dies sind Aktionen, die Sie normalerweise einmal bei der Anwendungsinitialisierung durchführen würden. Daher ist es nicht ganz fair, sie in das Timing aufzunehmen.

Es sollte auch beachtet werden, dass, wenn Sie diese Tests über Xcode mit der Metal Validation Layer und GPU Frame Capture aktiviert ausführen, das erhebliche Laufzeitkosten hat und die Ergebnisse zugunsten der CPU verfälschen wird.

Mit diesen Vorbehalten, hier ist, wie Sie posix_memalign verwenden können, um Speicher zuzuweisen, der verwendet werden kann, um eine MTLBuffer zu unterstützen. Der Trick besteht darin, sicherzustellen, dass der von Ihnen angeforderte Speicher tatsächlich page-aligned ist (dh seine Adresse ist ein Vielfaches von getpagesize()), was dazu führen kann, dass die Speichermenge größer ist als die tatsächliche Speicherkapazität Ihrer Daten:

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 
} 

Dies erfordert data eine UnsafeMutableBufferPointer<DataType> machen, anstatt eine [DataType], da Swift Array seine eigene Sicherungsspeicher zuordnet. Sie müssen auch die Anzahl der zu operierenden Datenelemente weitergeben, da die count des veränderbaren Pufferzeigers aufgerundet wurde, um den Puffer page-aligned zu machen.

Um einen MTLBuffer zu erstellen, der mit diesen Daten gesichert wird, verwenden Sie die API newBufferWithBytesNoCopy(_:length:options:deallocator:).Es ist wichtig, dass die von Ihnen angegebene Länge erneut ein Vielfaches der Seitengröße ist. andernfalls gibt diese Methode nil:

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

Hier stellen wir keine deallocator, aber Sie sollten den Speicher frei, wenn Sie es mit fertig sind, durch die baseAddress des Pufferzeiger auf free() vorbei.