iOS上の配列のスウィートメタル並列和計算

Question

@ Kametrixom answerに基づいて、配列内のsumの並列計算のためのテストアプリケーションをいくつか作成しました。

私のテストアプリケーションは、次のようになります。

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

そして、私はこのようになりますシェーダ持っている：私の驚きの機能sumParallel4で

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

を、私はそれがshouldnと思った、最速でありますそうです。私が関数sumParralelとsumParallel3を呼び出すと、関数の順序を変更しても、最初の関数は常に遅くなることに気付きました。 （sumParral3を呼び出すと、最初はこれが遅くなりますが、これは遅いです。）

これはなぜですか？なぜsumParallel3はsumParallelよりずっと高速ではありませんか？ sumParallel4はCPUで計算されますが、なぜ最速ですか？

---

posix_memalignでGPU機能を更新するにはどうすればよいですか？私はそれがGPUとCPU間で共有メモリを持つので、より速く動作するはずですが、私は魔女の配列をこのように（データまたは結果）割り当てなければならないと、データが関数内で渡されたパラメータであればposix_memalign ？

Answer 1

これらのテストをiPhone 6で実行すると、メタルバージョンはナイーブなCPU合計よりも3倍速いと2倍速いのが分かりました。私が以下に述べる変更を加えれば、それは一貫してより速くなりました。

メタルバージョンの実行にかかるコストは、バッファの割り当てだけでなく、デバイスの初回作成とパイプライン状態の計算にも起因すると考えられます。これらは、通常はアプリケーションの初期化時に一度実行するアクションなので、タイミングにそれらを含めることは完全に公正ではありません。

メタル検証レイヤーとGPUフレームキャプチャが有効になっているXcodeでこれらのテストを実行している場合、実行時に大きなコストがかかり、CPUに有利な結果が得られます。

これらの警告では、をバックアップするために使用できるメモリを割り当てるために、posix_memalignを使用する方法を次に示します。実際にリクエストするメモリがページ単位であることを確認することです（つまり、そのアドレスは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 
} 

これはスウィフトのArrayは独自のバッキングストアを割り当てているので、むしろ[DataType]より、dataUnsafeMutableBufferPointer<DataType>を作る必要がありません。変更可能なバッファポインタのcountがバッファをページアライメントするように切り上げられているため、操作するデータ項目の数を渡す必要もあります。

MTLBufferを実際に作成するには、newBufferWithBytesNoCopy(_:length:options:deallocator:) APIを使用してください。もう一度、あなたが提供する長さがページサイズの倍数であることが重要です。そうでない場合、このメソッドはnilを返します。

ここ

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

、我々はデアロケータを提供していませんが、あなたはそれを使用して行われているとき、あなたはfree()にバッファポインタのbaseAddressを渡すことによって、メモリを解放する必要があります。