【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作為默認值并返回該GPUdevice對象谓娃。在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庫中肘习。您可以使用MTLLibraryMTLFunction對象獲取有關(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),CUPGPU都可以訪問它。
要使用隨機數(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ū)⑵鋵懭胛募S捎谟嬎銉H用于說明創(chuàng)建Metal應用程序的過程忽媒,因此示例將讀取存儲在輸出緩沖區(qū)中的值并進行測試以確保CPUGPU計算出相同的結(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
  • 文/蒼蘭香墨 我猛地睜開眼癌蚁,長吁一口氣:“原來是場噩夢啊……” “哼!你這毒婦竟也來了兜畸?” 一聲冷哼從身側(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
  • 正文 年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