本來Cuda用的挺好，為了Apple，放棄Cuda，改投OpenCl。好不容易OpenCl也算熟悉了，WWDC2018又宣布了Metal2，建議大家放棄OpenCl，使用Metal Performance Shaders。
Apple是一個富有“革命性”創新力的公司，很多創新，會徹底的放棄原有的積累。不斷帶來新能力的同時，也讓人又愛又恨。

下面是一個例子，用於演示如何使用Metal+Shader來加速mac的大規模數據計算。
主程序使用swift。隨機生成一個大規模的整數數組，然後分配到GPU內核上並行對數組進行求和。
Metal部分的各項邏輯建議看官方文檔https://developer.apple.com/metal/，只重點說一下計算部分。計算是由Shader子程序（核函數）完成的，Shader編程所使用的語言衍生自c++14，所以跟cpu通訊所使用的數據結構基本都是使用c語言可以接受的類型。可以把Shader語言理解為c++的一個子集。官方的建議是可以有大量的計算，但盡力減少邏輯語句之類需要GPU核進行預判從而降低速度的內容。大多情況下單個內核的計算速度並不快，使用GPU加速計算的原因是GPU都具有很多個計算單元進行並行的計算。
通常在Shader函數的參數中，至少包含3個部分：輸入、輸出、進程的ID。前兩個參數好理解，第三個參數就是因為該核函數可能隨機的運行在某個GPU內核上進行計算工作，應當根據這個唯一的ID分配出來唯一的任務在程序中來計算，從而達到並發的效果。所以核函數都應當是支持並發、支持數據切割分塊計算。
Metal對並發的支持首先是線程組數量threadgroupsPerGrid，這個基本上是跟GPU核心數相關的，另一個是批次數量threadsPerThreadgroup，這個要求是線程組數量的整倍數。
其它的內容請看代碼中的註釋。主程序命名為testCompute.swift

import Metal

//定義數據集長度，總共count個數據求和
//swift數字立即數中可以添加下劃線表現出來科學計數法的方式，很有特色
let count = 10_000_000
//每elementsPerSum個數據分配到一個核匯總一次
let elementsPerSum = 10_000

//每個數據的類型，必須使用C兼容的類型，
//因為GPU運行的shader語言是從C++14衍生來的
typealias DataType = CInt

//設備，就是GPU
let device = MTLCreateSystemDefaultDevice()!
//載入當前目錄下的default.metallib(編譯後的shader),使用其中的parsum核函數
let parsum = device.makeDefaultLibrary()!.makeFunction(name:"parsum")!
//如果shader文件不是默認名稱，可以使用下面的方法載入指定文件
//let lib = try! dev.makeLibrary(filepath:"default.metallib")!.newFunctionWithName("parsum")!
//建立GPU運算的流水線
let pipeline = try! device.makeComputePipelineState(function:parsum)
//生成隨機數據集
var data = (0..<count).map{ _ in DataType(arc4random_uniform(100)) }
//傳遞給核函數的數據總數，所以也用C兼容方式
var dataCount = CUnsignedInt(count)
//傳遞給核函數的每組匯總數量，同上
var elementsPerSumC = CUnsignedInt(elementsPerSum)
//返回的分批匯總的結果數
let resultsCount = (count + elementsPerSum - 1) / elementsPerSum

//建立兩個同GPU通信的緩沖區，一個用於輸入給核函數，一個用用於核函數返回結果
let dataBuffer = device.makeBuffer(bytes:&data, length: MemoryLayout<CInt>.size * count, options: []) // Our data in a buffer (copied)
let resultsBuffer = device.makeBuffer(length:MemoryLayout<CInt>.size * resultsCount, options: []) // A buffer for individual results (zero initialized)
//返回結果是c指針，要轉換成swift可訪問的數據類型
let results = UnsafeBufferPointer<DataType>(
    start: resultsBuffer!.contents().assumingMemoryBound(to:CInt.self), count: resultsCount)
//建立GPU命令隊列
let queue = device.makeCommandQueue()
//GPU命令緩沖區，一般有多個運算會都放置到緩沖區，一次性提交執行
let cmds = queue!.makeCommandBuffer()
//命令編碼器是用於將一條GPU核函數調用的函數、參數等打包到一起
let encoder = cmds!.makeComputeCommandEncoder()!

//設置一條GPU核函數調用的函數及其相關參數，如前所述，必須使用C兼容的類型
encoder.setComputePipelineState(pipeline)
encoder.setBuffer(dataBuffer, offset: 0, index: 0)
encoder.setBytes(&dataCount, length: MemoryLayout<CUnsignedInt>.size, index: 1)
encoder.setBuffer(resultsBuffer, offset: 0, index: 2)
encoder.setBytes(&elementsPerSumC, length: MemoryLayout<CUnsignedInt>.size, index: 3)

//設定每組任務數量
let threadgroupsPerGrid = MTLSize(width: (resultsCount + pipeline.threadExecutionWidth - 1) / pipeline.threadExecutionWidth, height: 1, depth: 1)
//設定每批次任務數量，必須是上面組數量的整倍數
let threadsPerThreadgroup = MTLSize(width: pipeline.threadExecutionWidth, height: 1, depth: 1)
//分配任務線程
encoder.dispatchThreadgroups(threadgroupsPerGrid, threadsPerThreadgroup: threadsPerThreadgroup)
//完成一個調用的所有設置
encoder.endEncoding()

var start, end : UInt64
var result : DataType = 0

start = mach_absolute_time()
//真正提交任務
cmds!.commit()
//等待完成GPU計算
cmds!.waitUntilCompleted()
//GPU計算式分批次匯總的，數量已經很少了，最後用CPU進行完整的匯總
for elem in results {
    result += elem
}
end = mach_absolute_time()
//顯示GPU計算結果及所用時間
print("Metal result: \(result), time: \(Double(end - start) / Double(NSEC_PER_SEC))")
result = 0

//下面是使用CPU完整的計算一次，並顯示結果、耗費時間
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))")
 

shade程序命名為：shader.metal

//各項數據類型必須跟Swift中定義的相同
#include <metal_stdlib>
typedef unsigned int uint; 
typedef int DataType;

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;
    //計算本批次數據的開始結束位置
    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];
}

給一個在命令行使用的編譯腳本：

#!/bin/bash
xcrun metal -o shader.air shader.metal
xcrun metal-ar rcs shader.metal-ar shader.air
xcrun metallib -o default.metallib shader.metal-ar
swiftc testCompute.swift?   

在我的筆記本上運行效果如下:

metal> ./testCompute 
Metal result: 495056208, time: 0.017362745
CPU result: 495056208, time: 1.210801891

作為一個比較片面的比較，GPU計算速度，比CPU快121倍。
測試環境：
MacBook Pro (13-inch, 2017, Four Thunderbolt 3 Ports)
CPU:3.1 GHz Intel Core i5
Graphics:Intel Iris Plus Graphics 650 1536 MB
Memory:8 GB 2133 MHz LPDDR3
Xcode:9.4.1

參考資料：
https://stackoverflow.com/questions/38164634/compute-sum-of-array-values-in-parallel-with-metal-swift

Metal並行計算以及Metal程序的命令行編譯