【Metal】在GPU上執(zhí)行計算（I）

下載

概述

在此示例中绪励，您將學習所有Metal應用程序中使用的基本任務(wù)：

  • 將用C編寫簡單的函數(shù)轉(zhuǎn)換為Metal Shading Language(MSL)疏魏，以便它可以在GPU上運行
  • 找到GPU
  • 通過創(chuàng)建管道使MSL函數(shù)在GPU上運行
  • 創(chuàng)建GPU可訪問的內(nèi)存以保存數(shù)據(jù)
  • 創(chuàng)建命令緩沖區(qū)并編碼GPU命令以操縱數(shù)據(jù)
  • 將緩沖區(qū)提交到命令隊列以使GPU執(zhí)行編碼命令

編寫GPU函數(shù)以執(zhí)行計算

為了說明GPU編程，這個應用程序?qū)蓚€數(shù)組的想應元素添加到一起官份，將結(jié)果寫入第三個數(shù)組。清單1顯示了一個在CPU上執(zhí)行此計算的函數(shù)羔味，用C語言編寫忘蟹。它循環(huán)遍歷索引寒瓦，計算循環(huán)每次一迭代的結(jié)果杂腰。

清單1 數(shù)組添加，用C編寫

void add_arrays(const float* inA, 
                const float* inB,
                float* result,
                int length)
{
    for (int index = 0; index < length; index++) 
    { 
        result[index] = inA[index] + inB[index];
    }
}

每個值都是獨立計算的少辣，同時也可以安全地計算這些值。要在GPU上執(zhí)行計算忙干，您需要在Metal Shading Language(MSL)中重寫此功能。MSL是為GPU編程設(shè)計的C++的變體浪藻。在Metal中捐迫，在GPU上運行的代碼稱為著色器(shader)，因為歷史上他們首先用于計算3D圖形中的顏色爱葵。清單2顯示了MSL中的著色器施戴，他執(zhí)行與清單1相同的計算。示例項目add.metal文件中定義了此函數(shù)萌丈。Xcode構(gòu)建應用程序目標中的所有文件赞哗，并創(chuàng)建一個默認的Metal庫，它嵌入到您的應用程序中辆雾。您將在本實例的后面看到如何加載默認庫。

清單2 數(shù)組添加，用MSL編寫

kernel void add_arrays(device const float* inA,
                       device const float* inB,
                       device float* result,
                       uint index [[thread_position_in_grid]])
{
    // for循環(huán)被一組線程代替，每個線程都會調(diào)用此函數(shù)
    result[index] = inA[index] + inB[index];
}

清單1和清單2類似尖殃，但MSL版本有一些重要區(qū)別器躏。仔細看看清單2揽浙。

首先走敌，該函數(shù)添加一個kernel關(guān)鍵字，聲明該函數(shù)是：

  • 一個公共的GPU功能。公共函數(shù)是您的應用可以看到的唯一功能。其他著色器函數(shù)也無法調(diào)用公共函數(shù)展哭。
  • 一個計算函數(shù)（也稱為計算內(nèi)核），其執(zhí)行使用線程的網(wǎng)格中的并行計算手蝎。

參閱使用渲染管道渲染基元以了解用于聲明公共圖形函數(shù)的其他函數(shù)的關(guān)鍵字鞍时。

add_array函數(shù)使用device關(guān)鍵字聲明了三個參數(shù)笙僚，該關(guān)鍵字表示這些指針位于地址空間中。MSL為內(nèi)存定義了幾個不相交的地址空間。每當在MSL中聲明指針時锦援，都必須提供一個關(guān)鍵字來聲明其地址空間凉泄。使用device地址空間聲明GPU可以讀取和寫入持久內(nèi)存蒂誉。

清單2從清單1中刪除了for循環(huán)讥电，因為該函數(shù)現(xiàn)在將由計算網(wǎng)格中的多個線程調(diào)用月趟。此示例創(chuàng)建一個完全匹配數(shù)組維度的一維線程網(wǎng)絡(luò)。

要替換先前由for循環(huán)提供的索引看疗，該函數(shù)將index使用另一個MSL關(guān)鍵字thread_position_in_grid獲取一個新參數(shù)，該參數(shù)使用C++屬性語法指定是复。此關(guān)鍵字聲明Metal應為每一個線程計算唯一索引，并在此參數(shù)中傳遞該索引。因為add_arrays使用1D網(wǎng)格嘶居，索引被定義為標量整數(shù)炸客。即使刪除了循環(huán)，清單1和清單2也使用相同的代碼將兩個數(shù)字相加戈钢。如果要將類似的代碼從C或者C++轉(zhuǎn)化為MSL痹仙，請以相同的網(wǎng)格替換循環(huán)邏輯。

尋找一個GPU

在你的應用程序中殉了，MTLDevice對象是GPU的精簡抽象开仰，你用它來與GPU通信。Metal為每一個GPU創(chuàng)建了MTLDevice薪铜，你可以通過調(diào)用MTLCreateSystemDefaultDevice()獲取默認的device對象众弓。在macOS中，Mac可以有多個GPU隔箍，Metal可以選擇其中一個GPU作為默認值并返回該GPU的device對象谓娃。在macOS中，Metal提供了可用于檢索所有device對象的其他API鞍恢，但此示例僅適用默認值傻粘。

let device = MTLCreateSystemDefaultDevice()

初始化Metal對象

Metal將其他與GPU相關(guān)的實體（如編譯著色器(compiled shaders)、內(nèi)存緩沖區(qū)(memory buffers)和紋理(textures)）表示為對象帮掉。要創(chuàng)建這些特定于GPU的對象弦悉，可以調(diào)用MTLDevice的相關(guān)函數(shù)或者由MTLDevice創(chuàng)建的對象調(diào)用函數(shù)。由device對象直接或間接創(chuàng)建的所有對象僅可用于該device對象蟆炊。使用多個GPU的應用程序?qū)⑹褂枚鄠€device對象稽莉，并為每個device創(chuàng)建類似的Metal對象層次結(jié)構(gòu)。
示例應用程序使用自定義類MetalAdder來管理與GPU通信所需的對象涩搓。類的初始化程序創(chuàng)建這些對象并將它們存儲在其屬性中污秆。該應用程序創(chuàng)建此類的實例，傳入Metaldevice對象用以創(chuàng)建輔助對象昧甘。該MetalAdder對象保持對Metal對象的強引用良拼，直到它完成執(zhí)行。

let adder = MetalAdder(device: device)

在Metal中充边，昂貴的初始化任務(wù)可以運行一次庸推，保留結(jié)果并且使用成本低廉。您很少需要在性能敏感的代碼中運行此類任務(wù)浇冰。

獲取Metal功能的參考

初始化程序所做的第一件事是加載函數(shù)并準備它在GPU上運行贬媒。在構(gòu)建應用程序時，Xcode會編譯add_arrays函數(shù)并使其添加它嵌入應用程序的默認Metal庫中肘习。您可以使用MTLLibrary和MTLFunction對象獲取有關(guān)Metal庫及其中包含的函數(shù)信息际乘。要獲取add_arrays函數(shù)的對象，請要求MTLDevice來創(chuàng)建默認庫MTLLibrary對象漂佩，然后向庫請求表示著色器函數(shù)的MTLFunction對象脖含。

init?(device: MTLDevice) {
    super.init()
        
    mDevice = device
        
    guard let defaultLibrary = mDevice?.makeDefaultLibrary() else {
        print("Failed to find the default library.")
        return nil
    }
        
    guard let addFunction = defaultLibrary.makeFunction(name: "add_arrays") else {
        print("Failed to find adder function.")
        return nil
    }
}

準備Metal管道（Pipeline）

函數(shù)對象是MSL函數(shù)的代理罪塔，但它不是可執(zhí)行代碼。您可以通過創(chuàng)建管道(pipeline)將該函數(shù)轉(zhuǎn)為可執(zhí)行代碼器赞。管道指定GPU執(zhí)行以完成特定任務(wù)步驟垢袱。在Metal中，管道由管道狀態(tài)對象(pipeline state)表示港柜。由于此示例使用計算功能，因此應用程序會創(chuàng)建一個MTLComputePipelineState對象咳榜。

do {
    try mAddFunctionPOS = mDevice?.makeComputePipelineState(function: addFunction)
} catch {
    print("Failed to created pipline state object, error: \(error.localizedDescription).")
}

計算管道運行單個計算功能夏醉，可選地在運行函數(shù)之前操縱輸入數(shù)據(jù)，然后運行輸出數(shù)據(jù)涌韩。

創(chuàng)建管道狀態(tài)對象時畔柔，設(shè)備對象將完成為此特定的GPU*編譯的功能。此示例同步創(chuàng)建管道狀態(tài)對象并將其直接返回給應用程序臣樱。因為編譯確實需要一段時間靶擦，所以避免在性能敏感的代碼中同步創(chuàng)建管道狀態(tài)對象。

注意
Metal到目前為止您看到的代碼中返回的所有對象都將作為符合協(xié)議的對象返回雇毫。Metal使用協(xié)議來定義大多數(shù)特定于GPU的對象玄捕，以抽象到底層實現(xiàn)，這些類可能因不同的GPU而異棚放。Metal使用類定義與GPU無關(guān)的對象枚粘。任何給定的Metal協(xié)議的參考文檔都清楚的表明您是否可以在應用程序中實現(xiàn)該協(xié)議。

創(chuàng)建一個命令隊列（Command Queue）

要將工作發(fā)送到GPU飘蚯，您需要一個命令隊列馍迄。Metal使用隊列命令來安排命令。通過詢問一個MTLDevice來創(chuàng)建一個命令隊列局骤。

mCommandQueue = device.makeCommandQueue()

創(chuàng)建數(shù)據(jù)緩沖區(qū)和加載數(shù)據(jù)

初始化基本Metal對象后攀圈，加載要執(zhí)行的GPU數(shù)據(jù)。此任務(wù)的性能不太重要峦甩，但在應用程序啟動的早期仍然有用赘来。

GPU可以擁有自己的專用內(nèi)存，也可以與操作系統(tǒng)共享內(nèi)存穴店。Metal和操作系統(tǒng)內(nèi)核需要執(zhí)行額外的工作撕捍，以便將數(shù)據(jù)存儲在內(nèi)存中并使數(shù)據(jù)可供GPU使用。Metal使用資源對象抽象此內(nèi)存管理泣洞。（MTLResource）忧风。資源是GPU在運行命令時可以訪問的內(nèi)存分配。使用MTLDevice為其GPU創(chuàng)建資源球凰。

示例應用程序創(chuàng)建了三個緩存區(qū)狮腿，并使用隨機數(shù)據(jù)填充前兩個緩沖區(qū)腿宰。第三個緩沖將在add_arrays存儲其結(jié)果。

mBufferA = mDevice?.makeBuffer(length: bufferSize, options: .storageModeShared)
mBufferB = mDevice?.makeBuffer(length: bufferSize, options: .storageModeShared)
mBufferResult = mDevice?.makeBuffer(length: bufferSize, options: .storageModeShared)
        
generateRandomFloatData(buffer: mBufferA)
generateRandomFloatData(buffer: mBufferB)

此示例中的資源是（MTLBuffer）對象缘厢，她們是沒有預定義格式的內(nèi)存分配吃度。Metal將每個緩沖區(qū)管理為不透明的字節(jié)集合。但是贴硫，在著色器中使用緩沖區(qū)時指定格式椿每。這意味著您的著色器和您的應用需要就來回傳遞的任何數(shù)據(jù)的格式達成一致。
分配緩存區(qū)時英遭，您需要提供存儲模式以確定其某些性能特征以及CPU或者GPU是否可以訪問它间护。應用程序使用共享內(nèi)存（storageModeShared）,CUP和GPU都可以訪問它。
要使用隨機數(shù)據(jù)填充緩沖區(qū)挖诸，應用程序?qū)@取指向緩存區(qū)內(nèi)存的指針汁尺，并在CPU上將數(shù)據(jù)寫入其中。清單2中的add_arrays函數(shù)將參數(shù)聲明為浮點數(shù)的數(shù)組多律，因此您以相同的格式提供緩沖區(qū)：


- (void) generateRandomFloatData:(id<MTLBuffer>) buffer
{ 
    float* dataPtr = buffer.contents;
    for (unsigned long index = 0; index < arrayLength; index++)
    {
        dataPtr[index] = (float)rand()/(float)(RAND_MAX);
    }
}
/// Swift不會痴突，所以用OC

創(chuàng)建一個命令緩沖區(qū)

請求命令隊列創(chuàng)建命令緩沖區(qū)。

let commandBuffer = mCommandQueue?.makeCommandBuffer()

創(chuàng)建一個命令編碼器

要將命令寫入命令緩沖區(qū)狼荞，可以使用命令編碼器來處理要編碼的特定命令類型辽装。此示例創(chuàng)建一個計算命令編碼器，用于編碼計算傳遞粘秆。計算傳遞包含執(zhí)行計算管道的命令列表如迟。每個計算命令都會使GPU創(chuàng)建一個在GPU上執(zhí)行的線程網(wǎng)絡(luò)。

let commandEncoder = commandBuffer?.makeComputeCommandEncoder()

要對命令進行編碼攻走，請在編碼器上進行一系列方法調(diào)用殷勘。某些方法設(shè)置狀態(tài)信息，如管道狀態(tài)對象（PSO）或要傳遞給管道的參數(shù)昔搂。進行狀態(tài)更改后玲销，您將編碼命令以執(zhí)行管道。編碼器將所有狀態(tài)改變和命令參數(shù)寫入命令緩沖區(qū)摘符。

image

設(shè)置管道狀態(tài)和參數(shù)數(shù)據(jù)

設(shè)置要執(zhí)行命令的管道的管道狀態(tài)對象贤斜。然后為管道需要發(fā)送到add_arrays函數(shù)的任何參數(shù)設(shè)置數(shù)據(jù)。對于此管道逛裤，這意味著提供對三個緩沖區(qū)的引用瘩绒。Metal以參數(shù)出現(xiàn)在清單2中的函數(shù)聲明中的順序自動為緩沖區(qū)參數(shù)指定索引，從0開始带族，您使用相同的索引提供參數(shù)锁荔。

commandEncoder.setComputePipelineState(mAddFunctionPOS!)
commandEncoder.setBuffer(mBufferA, offset: 0, index: 0)
commandEncoder.setBuffer(mBufferB, offset: 0, index: 1)
commandEncoder.setBuffer(mBufferResult, offset: 0, index: 2)

您還可以為每個參數(shù)指定偏移量。偏移0意味著命令將從緩沖區(qū)來存儲多個參數(shù)蝙砌，為某個參數(shù)指定偏移量阳堕。
您沒有為index參數(shù)指定任何數(shù)據(jù)跋理，因為add_arrays函數(shù)將其值定義為由GPU提供。

指定線程計數(shù)和組織

接下來恬总，確定要創(chuàng)建的線程數(shù)以及如何組織這些線程前普。Metal可以創(chuàng)建1D、2D和3D網(wǎng)絡(luò)壹堰。add_array函數(shù)使用一維數(shù)組(dataSize x 1 x 1)拭卿，Metal從該網(wǎng)格生成0到-1之間的索引。

指定線程組大小

Metal將網(wǎng)格細分為稱為線程組的較小網(wǎng)格缀旁。Metal可以將線程組分派到CPU上的不同處理元素记劈，以加快處理速度。您還需要確定為命令創(chuàng)建線程組的大小并巍。

let gridSize = MTLSize(width: arrayLenght, height: 1, depth: 1)
var threadGroupSizeInt = mAddFunctionPOS?.maxTotalThreadsPerThreadgroup
assert(threadGroupSizeInt != nil, "Failed to get thread group size.")
        
if threadGroupSizeInt! > arrayLenght {
threadGroupSizeInt = arrayLenght
}
let threadGroupSize = MTLSize(width: arrayLenght, height: 1, depth: 1)

應用程序管道狀態(tài)對象請求最大可能的線程組，并在該大小大于數(shù)據(jù)集時收縮它换途。該maxTotalThreadsPerThreadgroup屬性給出線程組允許的最大線程數(shù)懊渡，這取決于用于創(chuàng)建管道狀態(tài)對象的函數(shù)的復雜性。

編寫Compute命令以執(zhí)行線程

最后军拟，編碼命令以調(diào)度線程網(wǎng)絡(luò)剃执。

commandEncoder.dispatchThreads(gridSize, threadsPerThreadgroup: threadGroupSize)

當GPU執(zhí)行此命令時，它使用您先前設(shè)置的狀態(tài)和命令的參數(shù)來分派線程以執(zhí)行計算懈息。

您可以使用編碼器執(zhí)行相同的步驟肾档，將多個計算命令編碼到計算傳遞中，而不執(zhí)行任何冗余步驟辫继。例如怒见，您可以設(shè)置一次管道狀態(tài)對象，然后為要處理的每個緩沖區(qū)集合設(shè)置參數(shù)并編碼命令姑宽。

結(jié)束計算通行證

如果沒有其他命令添加到計算傳遞遣耍，則結(jié)束編碼過程以關(guān)閉計算傳遞。

commandEncoder?.endEncoding()

提交命令緩沖區(qū)以執(zhí)行命令

通過將命令緩沖區(qū)提交到隊列來運行命令緩沖區(qū)中的命令炮车。

commandBuffer?.commit()

命令隊列創(chuàng)建了命令緩沖區(qū)舵变，因此提交緩沖區(qū)始終將其放在該隊列上。提交命令緩沖區(qū)后瘦穆，Metal異步準備執(zhí)行命令纪隙，然后調(diào)度命令緩沖區(qū)以在GPU上執(zhí)行。GPU執(zhí)行命令緩沖區(qū)中的所有命令后扛或，Metal將命令緩沖區(qū)標記為完成绵咱。

等待計算完成

當GPU處理您的命令時，您的應用程序可以執(zhí)行其他工作告喊。此示例不需要執(zhí)行任何其他工作麸拄，因此只需等待命令緩沖區(qū)完成派昧。

commandBuffer?.waitUntilCompleted()

或者，要在Metal處理完所有命令時收到通知拢切，請在命令緩沖區(qū)中添加完成處理程序(addCompletedHandler(_:))蒂萎，或者通過讀取其status來檢查命令緩沖區(qū)的狀態(tài)。

從緩沖區(qū)讀取狀態(tài)

命令緩沖區(qū)完成后淮椰，GPU的計算存儲在輸出緩沖區(qū)中五慈，Metal執(zhí)行任何必要的步驟以確保CPU可以看到他們。在真實的應用程序中主穗，您將從緩沖區(qū)讀取結(jié)果并對其執(zhí)行某些操作泻拦，例如在屏幕上顯示結(jié)果或?qū)⑵鋵懭胛募Ｓ捎谟嬎銉H用于說明創(chuàng)建Metal應用程序的過程忽媒，因此示例將讀取存儲在輸出緩沖區(qū)中的值并進行測試以確保CPU和GPU計算出相同的結(jié)果争拐。

- (void) verifyResults
{
    float* a = mBufferA.contents;
    float* b = mBufferB.contents;
    float* result = mBufferResult.contents;
    for (unsigned long index = 0; index < arrayLength; index++) 
    {
        assert(result[index] = a[index] + b[index]);
    }
}
/// Swift不會，所以用OC
最后編輯于 ：
?著作權(quán)歸作者所有,轉(zhuǎn)載或內(nèi)容合作請聯(lián)系作者
  • 人面猴
    序言：七十年代末晦雨，一起剝皮案震驚了整個濱河市架曹，隨后出現(xiàn)的幾起案子，更是在濱河造成了極大的恐慌闹瞧，老刑警劉巖绑雄，帶你破解...
    沈念sama閱讀 222,681評論 6贊 517
  • 死咒
    序言：濱河連續(xù)發(fā)生了三起死亡事件，死亡現(xiàn)場離奇詭異奥邮，居然都是意外死亡万牺，警方通過查閱死者的電腦和手機，發(fā)現(xiàn)死者居然都...
    沈念sama閱讀 95,205評論 3贊 399
  • 救了他兩次的神仙讓他今天三更去死
    文/潘曉璐 我一進店門洽腺，熙熙樓的掌柜王于貴愁眉苦臉地迎上來脚粟，“玉大人，你說我怎么就攤上這事已脓∩郝ィ” “怎么了？”我有些...
    開封第一講書人閱讀 169,421評論 0贊 362
  • 道士緝兇錄：失蹤的賣姜人
    文/不壞的土叔 我叫張陵度液，是天一觀的道長厕宗。 經(jīng)常有香客問我，道長堕担，這世上最難降的妖魔是什么已慢？ 我笑而不...
    開封第一講書人閱讀 60,114評論 1贊 300
  • ?港島之戀（遺憾婚禮）
    正文 為了忘掉前任，我火速辦了婚禮霹购，結(jié)果婚禮上佑惠，老公的妹妹穿的比我還像新娘。我一直安慰自己，他們只是感情好膜楷，可當我...
    茶點故事閱讀 69,116評論 6贊 398
  • 惡毒庶女頂嫁案：這布局不是一般人想出來的
    文/花漫 我一把揭開白布旭咽。 她就那樣靜靜地躺著，像睡著了一般赌厅。 火紅的嫁衣襯著肌膚如雪穷绵。 梳的紋絲不亂的頭發(fā)上，一...
    開封第一講書人閱讀 52,713評論 1贊 312
  • 城市分裂傳說
    那天特愿，我揣著相機與錄音仲墨，去河邊找鬼。 笑死揍障，一個胖子當著我的面吹牛目养，可吹牛的內(nèi)容都是我干的。 我是一名探鬼主播毒嫡，決...
    沈念sama閱讀 41,170評論 3贊 422
  • 雙鴛鴦連環(huán)套：你想象不到人心有多黑
    文/蒼蘭香墨 我猛地睜開眼癌蚁，長吁一口氣：“原來是場噩夢啊……” “哼！你這毒婦竟也來了兜畸？” 一聲冷哼從身側(cè)響起匈勋，我...
    開封第一講書人閱讀 40,116評論 0贊 277
  • 萬榮殺人案實錄
    序言：老撾萬榮一對情侶失蹤，失蹤者是張志新（化名）和其女友劉穎膳叨，沒想到半個月后，有當?shù)厝嗽跇淞掷锇l(fā)現(xiàn)了一具尸體痘系，經(jīng)...
    沈念sama閱讀 46,651評論 1贊 320
  • ?護林員之死
    正文 獨居荒郊野嶺守林人離奇死亡菲嘴，尸身上長有42處帶血的膿包…… 初始之章·張勛 以下內(nèi)容為張勛視角 年9月15日...
    茶點故事閱讀 38,714評論 3贊 342
  • ?白月光啟示錄
    正文 我和宋清朗相戀三年，在試婚紗的時候發(fā)現(xiàn)自己被綠了汰翠。 大學時的朋友給我發(fā)了我未婚夫和他白月光在一起吃飯的照片龄坪。...
    茶點故事閱讀 40,865評論 1贊 353
  • 活死人
    序言：一個原本活蹦亂跳的男人離奇死亡，死狀恐怖复唤，靈堂內(nèi)的尸體忽然破棺而出健田，到底是詐尸還是另有隱情，我是刑警寧澤佛纫，帶...
    沈念sama閱讀 36,527評論 5贊 351
  • ?日本核電站爆炸內(nèi)幕
    正文 年R本政府宣布妓局，位于F島的核電站，受9級特大地震影響呈宇，放射性物質(zhì)發(fā)生泄漏好爬。R本人自食惡果不足惜，卻給世界環(huán)境...
    茶點故事閱讀 42,211評論 3贊 336
  • 男人毒藥：我在死后第九天來索命
    文/蒙蒙 一甥啄、第九天 我趴在偏房一處隱蔽的房頂上張望存炮。 院中可真熱鬧，春花似錦、人聲如沸穆桂。這莊子的主人今日做“春日...
    開封第一講書人閱讀 32,699評論 0贊 25
  • 一樁弒父案，背后竟有這般陰謀
    文/蒼蘭香墨 我抬頭看了看天上的太陽享完。三九已至灼芭，卻和暖如春，著一層夾襖步出監(jiān)牢的瞬間驼侠，已是汗流浹背姿鸿。 一陣腳步聲響...
    開封第一講書人閱讀 33,814評論 1贊 274
  • 情欲美人皮
    我被黑心中介騙來泰國打工， 沒想到剛下飛機就差點兒被人妖公主榨干…… 1. 我叫王不留倒源，地道東北人苛预。 一個月前我還...
    沈念sama閱讀 49,299評論 3贊 379
  • 代替公主和親
    正文 我出身青樓，卻偏偏與公主長得像笋熬，于是被迫代替她去往敵國和親热某。 傳聞我的和親對象是個殘疾皇子，可洞房花燭夜當晚...
    茶點故事閱讀 45,870評論 2贊 361