[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中也稱為wavefrontwarp)是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 章完成。

下次見!

相關文章