[MetalKit]46-Introduction to compute using Metal 用 Metal 進行計算的簡介

蘋果API搬運工發表於2019-03-24

本系列文章是對 metalkit.org 上面MetalKit內容的全面翻譯和學習.

MetalKit系統文章目錄

在 GPU 程式設計的領域中，計算或者說GPGPU，是 GPU 程式設計中除渲染外的另一種用途。它們都涉及到了 GPU 並行程式設計，不同之處在於在計算中對執行緒的工作方式進行了更精細的控制。這樣，當你想要某些執行緒來處理問題的某一部分，同時其他執行緒去處理該問題的另一部分時，就會很有用。

本文是一系列關於計算的文章的開始篇。本文中的主題是關於影象處理，因為它是引入計算和執行緒管理的最簡單方法。

注意：本文假設您知道如何建立一個微型的Metal專案或playground，可以將螢幕清除為純色。

第一個不同點就是，你需要建立一個MTLComputePipelineState以取代以前渲染時用的MTLRenderPipelineState：

let function = library.makeFunction(name: "compute")
let pipelineState = device.makeComputePipelineState(function: function)
複製程式碼

第二件事是，你需要一個紋理，以供執行緒使用。如果你使用的是playground，那你只需要下面幾行：

let textureLoader = MTKTextureLoader(device: device)
let url = Bundle.main.url(forResource: "nature", withExtension: "jpg")!
let image = try textureLoader.newTexture(URL: url, options: [:])
複製程式碼

第三件事，你需要一個MTLComputeCommandEncoder物件，以便將先前建立的管線狀態物件和紋理，都附著上去：

commandEncoder.setComputePipelineState(pipelineState)
commandEncoder.setTexture(image, index: 0)
複製程式碼

第四件事，你需要一個kernel shader核心著色器，要記得，你之前開始時就為其建立了一個名為compute的函式。當然，你可以將核心程式碼放到 .metal檔案裡：

kernel void compute(texture2d<float, access::read> input [[texture(0)]],
                    texture2d<float, access::write> output [[texture(1)]],
                    uint2 id [[thread_position_in_grid]]) {
    float4 color = input.read(id);
    output.write(color, id);
}
複製程式碼

在著色程式碼中，輸入是你先前建立的MTLTexture物件，稱為image，輸出是一個可繪製紋理，你將向其中寫入資料，然後就可以被呈現到螢幕上了：

let drawable = view.currentDrawable
commandEncoder.setTexture(drawable.texture, index: 1)
複製程式碼

第五件事也是最後一件事是，你需要排程執行緒來幹活。有趣的事情就從現在開始了！你需要做的是在commandEncoder中結束編碼之前，加上幾句程式碼：

let threadsPerGroup = MTLSizeMake(100, 10, 1)
let groupsPerGrid = MTLSizeMake(15, 90, 1)
commandEncoder.dispatchThreadgroups(groupsPerGrid, threadsPerThreadgroup: threadsPerGroup)
複製程式碼

那麼這裡是怎麼做的呢？執行緒是以網格(grid)形式來排程處理資料的，網格可以是 1-，2-，或3-維的。在本例中，你用的是 2D 的網格，因為要處理的是一張圖片。不考慮維度的話，網格總是分割成多個執行緒組的，如下面的公式：

gridSize = groupsPerGrid * threadsPerGroup
複製程式碼

在本例中，你定義一個組有100 x10個執行緒，每個網格有15 x 90組。如果你執行你的 playground，你會看到類似下面的情況：

[MetalKit]46-Introduction to compute using Metal 用 Metal 進行計算的簡介

邊上的紅色是什麼東西？這是因為你試圖去猜測圖片的尺寸大小而導致的問題，執行緒數和組數應該用更“聰明”的方式獲取。

顯然，影象在兩個維度上都大於分派的執行緒數。您可以做的一件事是使用影象大小進行有根據的猜測，以獲得真正應該使用的組數量：

let width = Int(view.drawableSize.width)
let height = Int(view.drawableSize.height)
let w = threadsPerGroup.width
let h = threadsPerGroup.height
let groupsPerGrid = MTLSizeMake(width / w, height / h, 1)
複製程式碼

執行一下，圖片看起來會好很多了：

[MetalKit]46-Introduction to compute using Metal 用 Metal 進行計算的簡介

這裡又出現一個新的問題---利用不足。請看下圖的圖表：

[MetalKit]46-Introduction to compute using Metal 用 Metal 進行計算的簡介

通常，您會認為正確設計的網格是3 x 2組，每組4 x 4個執行緒，因此網格為12 x 8個執行緒。然而，底部和右側邊緣的一些螺紋未得到充分利用，因為它們沒有工作要做。

如果你製作一個較小的網格，比如8 x 4，它將會填滿整個組，又會產生你在開始時看到的紅色條帶。這意味著唯一可接受的解決方案是修復未充分利用問題。您可以通過在每個維度中新增額外的組來解決此問題，如下所示：

let groupsPerGrid = MTLSizeMake((width + w - 1) / w, (height + h - 1) / h, 1) 
複製程式碼

你所做的就是用(w-1, h-1, 1)來實際擴大網格尺寸。這又帶來了另一個風險 --- 訪問越界座標。要處理這個問題，您需要在讀取輸入影象之前向核心著色器新增邊界檢查：

if (id.x >= output.get_width() || id.y >= output.get_height()) {
    return;
}
複製程式碼

這將處理那些不應該做任何工作的執行緒，並處理越界的訪問。

那個執行緒組的大小怎麼樣 --- 無法優化嗎？到目前為止，你一直在猜這些尺寸。當然，還有一種方法可以獲得最佳的群組尺寸。硬體提供了一些可以通過管道狀態物件(pipeline state object)訪問的功能：

var w = pipelineState.threadExecutionWidth
var h = pipelineState.maxTotalThreadsPerThreadgroup / w
let threadsPerGroup = MTLSizeMake(w, h, 1)
複製程式碼

執行緒執行寬度（在其他API中也稱為wavefront或warp）是GPU組合在一起的執行緒數，因此它們可以並行地在不同的資料上執行相同的指令。組中的執行緒數應該是threadExecutionWidth的倍數，但絕不能大於maxTotalThreadsPerThreadgroup。

那太棒了！如何找到辦法，來避免做這些未充分利用和邊界檢查呢？Metal 也在這裡給你提供了幫助。 無需使用dispatchThreadgroups()，API提供了更新的dispatchThreads()函式，它實現了兩件事：

  1. 通過自動建立非均勻執行緒組（例如3 x 4）來適應邊緣情況，這樣就避免讓你處理未充分利用的問題。
  2. 它甚至可以決定需要多少組，前提是您為其提供網格大小和您想要使用的組大小。

注意：dispatchThreads()函式適用於所有macOS裝置，但它不適用於使用A10或更舊處理器的iOS裝置。

你需要做的就是，就下面程式碼替換計算每個網格組數的程式碼：

w = Int(view.drawableSize.width)
h = Int(view.drawableSize.height)
let threadsPerGrid = MTLSizeMake(w, h, 1)
commandEncoder.dispatchThreads(threadsPerGrid, threadsPerThreadgroup: threadsPerGroup)

複製程式碼

但是等一下，我是不是說過：這裡是最好玩的地方？是的，然後來到 kernel shader 中，移除邊界檢查程式碼，因為現在已經不需要它了。然後在最後一行前，新增下面程式碼，倒轉顏色通道：

color = float4(color.g, color.b, color.r, 1.0);
複製程式碼

執行一下 playground，你會看到類似下面的影象：

[MetalKit]46-Introduction to compute using Metal 用 Metal 進行計算的簡介

將上一行用下面程式碼替換，它將灰度應用於影象：

color.xyz = (color.r * 0.3 + color.g * 0.6 + color.b * 0.1) * 1.5;
複製程式碼

執行一下 playground，你會看到類似下面的影象：

[MetalKit]46-Introduction to compute using Metal 用 Metal 進行計算的簡介

最後，將下面程式碼替換：

float4 color = input.read(id);
color.xyz = (color.r * 0.3 + color.g * 0.6 + color.b * 0.1) * 1.5;
複製程式碼

替換為下面的程式碼，這裡將圖片將影象畫素化為5畫素的正方形：

uint2 index = uint2((id.x / 5) * 5, (id.y / 5) * 5);
float4 color = input.read(index);
複製程式碼

執行一下 playground，你會看到類似下面的影象：

[MetalKit]46-Introduction to compute using Metal 用 Metal 進行計算的簡介

玩得開心麼？希望你玩得開心。如果你想要學習更多關於影象處理的知識，Simon Gladman有一本好書，Core Image For Swift。本文只是一個對 GPGPU 和GPU計算功能的簡短介紹。請繼續關注新主題。

原始碼已經發布在Github上。本文基於書籍Metal by Tutorials的第 16 章完成。

下次見！

相關文章