Metal: 使用圖形處理器渲染高級(jí)三維圖形并執(zhí)行數(shù)據(jù)并行計(jì)算
編寫(xiě)一個(gè)GPU函數(shù)來(lái)執(zhí)行運(yùn)算砂心。
為了說(shuō)明GPU編程，此應(yīng)用程序?qū)蓚€(gè)數(shù)組的相應(yīng)元素添加到一起辑甜。把結(jié)果記錄在第三個(gè)數(shù)組里面，在列表1一種顯示一個(gè)用C語(yǔ)言寫(xiě)的靶端， 運(yùn)行在在CPU上的函數(shù)送爸。它通過(guò)索引循環(huán)，每次循環(huán)迭代計(jì)算一個(gè)值。
列表1 數(shù)組元素相加抚官，C語(yǔ)言編寫(xiě)
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];
}
}
每個(gè)值都是獨(dú)立計(jì)算的，所以每個(gè)值會(huì)同時(shí)進(jìn)行安全的計(jì)算阶捆。
為了讓運(yùn)算執(zhí)行在GPU上凌节，需要通過(guò)MSL（Metal Shading Language），MSL是為GPU編程而設(shè)計(jì)的C++變體洒试，在Metal API 中在GPU上運(yùn)行的代碼被稱(chēng)為著色器倍奢，因?yàn)樵跉v史上他們是第一次被用來(lái)運(yùn)算3D繪圖上顏色。在表2中顯示MSL上的一個(gè)執(zhí)行和表1一樣的運(yùn)算的著色器垒棋。
列表2卒煞，數(shù)組的相加荔燎，SML編寫(xiě)
kernel void add_arrays(device const float* inA,
device const float* inB,
device float* result,
uint index [[thread_position_in_grid]])
{
// the for-loop is replaced with a collection of threads, each of which
// calls this function.
result[index] = inA[index] + inB[index];
}
列表1和列表2 是相似的再层，但是在SML版本中有一些重要的不同點(diǎn)。我們來(lái)看一下列表2.
首先酿联，函數(shù)添加kernel關(guān)鍵字乖订，該關(guān)鍵字聲明函數(shù)為：
1.一個(gè)公共GPU函數(shù)扮饶，公共函數(shù)是在APP上能看見(jiàn)的唯一的函數(shù)， 公共函數(shù)也不能被其他的著色器函數(shù)調(diào)用乍构。
2.一個(gè)運(yùn)算函數(shù)（也被稱(chēng)之為運(yùn)算內(nèi)核）甜无，是通過(guò)線程網(wǎng)絡(luò)進(jìn)行平行運(yùn)算的函數(shù)。
閱讀使用渲染管道渲染基本體去學(xué)習(xí)其他更多的公共繪圖函數(shù)的函數(shù)關(guān)鍵字哥遮。
add_arrays函數(shù)通過(guò)設(shè)備關(guān)鍵字聲明了它的三個(gè)屬性毫蚓，這說(shuō)明這些屬性的指針在設(shè)備的地址空間里。MSL為內(nèi)存定義了幾個(gè)不連續(xù)的地址空間昔善。無(wú)論何時(shí)， 你在MSL上定義一個(gè)指針畔乙，你必須提供一個(gè)關(guān)鍵字來(lái)定義它的地址空間君仆。使用設(shè)備地址空間聲明GPU可以讀寫(xiě)的持久內(nèi)存。
列表2 移除了列表1中的for循環(huán)牲距，因?yàn)檫@個(gè)函數(shù)現(xiàn)在在運(yùn)算網(wǎng)絡(luò)中的多個(gè)線程調(diào)用返咱，此示例創(chuàng)建與數(shù)組維度完全匹配的線程的一維網(wǎng)格，因此數(shù)組中的每個(gè)條目都由不同的線程計(jì)算牍鞠。為了替換for循環(huán)提供的上一個(gè)索引咖摹，通過(guò)使用C++屬性語(yǔ)法置頂?shù)牧硪粋€(gè)MSL關(guān)鍵字（thread_position_in_grid）給函數(shù)設(shè)置一個(gè)新的索引屬性，此關(guān)鍵字聲明Metal 應(yīng)該為每一個(gè)線程計(jì)算一個(gè)唯一的索引难述，并在該參數(shù)中傳遞該索引萤晴。因?yàn)閍dd_arrays用的是一維網(wǎng)格吐句，這個(gè)索引會(huì)被定義為標(biāo)量整數(shù)，即便循環(huán)被移除店读，清單1和清單2使用同一行代碼將這兩個(gè)數(shù)字相加嗦枢。如果你想從C和C++中轉(zhuǎn)換一樣的代碼到SML，用同樣的方法把循環(huán)邏輯替換為網(wǎng)格屯断。
在應(yīng)用程序中文虏，MTLDevice對(duì)象是GPU的一個(gè)抽象,用它可以和GPU交互，Metal 為每一個(gè)GPU創(chuàng)建一個(gè)MTLDevice殖演，你可以調(diào)用 MTLCreateSystemDefaultDevice()獲得默認(rèn)device 對(duì)象氧秘。在macOS 中，Mac 可以有多個(gè)GPU趴久，Metal 選擇其中一個(gè)作為默認(rèn)GPU并返回他的device 對(duì)象丸相。在macOS中，Metal還提供了其他的APIs用來(lái)檢索所有的device 對(duì)象朋鞍，但下面的例子只用于默認(rèn)的對(duì)象已添。
id<MTLDevice> device = MTLCreateSystemDefaultDevice();
初始化Metal對(duì)象
初始化做的第一件事兒是加載函數(shù)并準(zhǔn)備它在GPU 上運(yùn)行，當(dāng)你創(chuàng)建應(yīng)用滥酥，Xcode 編譯add_arrays函數(shù)并且把它添加到嵌入到應(yīng)用里面的metal庫(kù)中更舞。你可以用MTLLibrary 和 MTLFunction 對(duì)象來(lái)獲取metal庫(kù)和函數(shù)中的相關(guān)信息。為了獲取一個(gè)表示add_arrays的函數(shù)坎吻，通過(guò) MTLDevice為默認(rèn)庫(kù)創(chuàng)建一個(gè)MTLLibrary對(duì)象并向庫(kù)請(qǐng)求一個(gè)MTLFunction對(duì)象來(lái)表示著色器函數(shù)缆蝉。
- (instancetype) initWithDevice: (id<MTLDevice>) device
{
self = [super init];
if (self)
{
_mDevice = device;
NSError* error = nil;
// Load the shader files with a .metal file extension in the project
id<MTLLibrary> defaultLibrary = [_mDevice newDefaultLibrary];
if (defaultLibrary == nil)
{
NSLog(@"Failed to find the default library.");
return nil;
}
id<MTLFunction> addFunction = [defaultLibrary newFunctionWithName:@"add_arrays"];
if (addFunction == nil)
{
NSLog(@"Failed to find the adder function.");
return nil;
}
準(zhǔn)備一個(gè)metal管道
函數(shù)對(duì)象是MSL函數(shù)的一個(gè)代理，但是它不是一個(gè)可執(zhí)行的代碼瘦真， 通過(guò)創(chuàng)建一個(gè)管道把函數(shù)轉(zhuǎn)換成可執(zhí)行代碼刊头。管道指定了GPU完成特定任務(wù)所執(zhí)行的步驟。在Metal 中诸尽，管道由管道狀態(tài)對(duì)象表示原杂。因?yàn)檫@個(gè)例子用了一個(gè)計(jì)算函數(shù)，應(yīng)用創(chuàng)建一個(gè)MTLComputePipelineState對(duì)象您机。
_mAddFunctionPSO = [_mDevice newComputePipelineStateWithFunction: addFunction error:&error];
一個(gè)計(jì)算管道運(yùn)行一個(gè)簡(jiǎn)單的計(jì)算函數(shù)穿肄，在運(yùn)行函數(shù)之前可選地操作輸入數(shù)據(jù)，以及在運(yùn)行函數(shù)之后操作輸出數(shù)據(jù)际看。
當(dāng)創(chuàng)建一個(gè)管道狀態(tài)對(duì)象咸产，device對(duì)象完成這個(gè)特定GPU的函數(shù)編譯。此示例同步創(chuàng)建管道狀態(tài)對(duì)象仲闽，并將其直接返回給應(yīng)用程序脑溢。因?yàn)榫幾g確實(shí)需要一段時(shí)間，所以避免在對(duì)性能敏感的代碼中同步創(chuàng)建管道狀態(tài)對(duì)象赖欣。
提示
到目前為止您看到的代碼中Metal返回的所有對(duì)象都是作為符合協(xié)議的對(duì)象返回的屑彻。Metal使用協(xié)議來(lái)抽象底層實(shí)現(xiàn)類(lèi)來(lái)定義大多數(shù)特定于gpu的對(duì)象验庙，這些實(shí)現(xiàn)類(lèi)對(duì)于不同的gpu可能會(huì)有所不同。Metal使用類(lèi)定義了與gpu無(wú)關(guān)的對(duì)象酱酬，任何給定的Metal協(xié)議的參考文檔都明確說(shuō)明了您是否可以在應(yīng)用程序中實(shí)現(xiàn)該協(xié)議壶谒。
創(chuàng)建命令隊(duì)列
為了把工作發(fā)送到GPU，你需要?jiǎng)?chuàng)建一個(gè)命令隊(duì)列膳沽。metal使用命令隊(duì)列來(lái)調(diào)度命令汗菜，通過(guò)向MTLDevice請(qǐng)求一個(gè)命令隊(duì)列來(lái)創(chuàng)建一個(gè)命令隊(duì)列。
_mCommandQueue = [_mDevice newCommandQueue];
創(chuàng)建數(shù)據(jù)緩沖區(qū)并加載數(shù)據(jù)
初始化基本metal 對(duì)象之后挑社，為GPU 加載執(zhí)行的數(shù)據(jù)陨界，這個(gè)任務(wù)對(duì)性能的影響較小，但在應(yīng)用啟動(dòng)的早期這樣做仍然很有用痛阻。
GPU有它自己的專(zhuān)用內(nèi)存菌瘪，或者它可以與操作系統(tǒng)共享內(nèi)存。Metal和操作系統(tǒng)內(nèi)核需要執(zhí)行額外的工作阱当，以便讓您在內(nèi)存中存儲(chǔ)數(shù)據(jù)俏扩，并使這些數(shù)據(jù)可供GPU使用。Metal使用資源對(duì)象抽象了這個(gè)內(nèi)存管理弊添。MTLResource是GPU在運(yùn)行命令時(shí)可以訪問(wèn)的內(nèi)存分配的資源录淡。通過(guò) MTLDevice來(lái)為GPU創(chuàng)建資源。
示例應(yīng)用創(chuàng)建三個(gè)緩沖區(qū)并用隨機(jī)數(shù)據(jù)填充前兩個(gè)油坝，第三個(gè)緩沖區(qū)是add_arrays存儲(chǔ)結(jié)果的地方嫉戚。
_mBufferA = [_mDevice newBufferWithLength:bufferSize options:MTLResourceStorageModeShared];
_mBufferB = [_mDevice newBufferWithLength:bufferSize options:MTLResourceStorageModeShared];
_mBufferResult = [_mDevice newBufferWithLength:bufferSize options:MTLResourceStorageModeShared];
[self generateRandomFloatData:_mBufferA];
[self generateRandomFloatData:_mBufferB];
這個(gè)示例中的資源是MTLBuffer對(duì)象，它是沒(méi)有預(yù)定義格式的內(nèi)存分配澈圈。Metal將每個(gè)緩沖區(qū)作為一個(gè)不透明的字節(jié)集合來(lái)管理彬檀。但是，在著色器中使用緩沖區(qū)時(shí)你來(lái)指定他的格式瞬女。這意味著你的著色器和你的應(yīng)用程序需要就任何來(lái)回傳遞的數(shù)據(jù)的格式達(dá)成一致窍帝。
當(dāng)你分配一個(gè)緩沖區(qū)，你提供一個(gè)存儲(chǔ)模式來(lái)決定它的一些性能特征以及CPU或GPU是否可以訪問(wèn)它诽偷。示例應(yīng)用程序使用共享內(nèi)存storageModeShared坤学，CPU和GPU都可以訪問(wèn)該內(nèi)存。
為了用隨機(jī)數(shù)據(jù)填充緩沖區(qū)渤刃，應(yīng)用程序從緩沖區(qū)內(nèi)存獲取一個(gè)指針并在CPU上給它寫(xiě)入數(shù)據(jù)。在列表2中的add_array函數(shù)聲明了它的浮點(diǎn)數(shù)類(lèi)型的數(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);
}
}
創(chuàng)建一個(gè)命令緩沖區(qū)
請(qǐng)求命令隊(duì)列創(chuàng)建一個(gè)命令緩沖區(qū)。
id<MTLCommandBuffer> commandBuffer = [_mCommandQueue commandBuffer];
創(chuàng)建一個(gè)命令編碼器
為了給命令緩沖區(qū)寫(xiě)入命令刑峡，你可以對(duì)要編碼的特定類(lèi)型的命令使用命令編碼器洋闽。這個(gè)示例創(chuàng)建一個(gè)計(jì)算命令編碼器玄柠，用來(lái)編碼計(jì)算過(guò)程。每個(gè)計(jì)算命令都會(huì)導(dǎo)致GPU創(chuàng)建線程網(wǎng)格以在GPU上執(zhí)行诫舅。
id<MTLComputeCommandEncoder> computeEncoder = [commandBuffer computeCommandEncoder];
為了編碼一個(gè)命令羽利，對(duì)編碼器進(jìn)行一系列方法調(diào)用。有些方法設(shè)置狀態(tài)信息刊懈，如管道狀態(tài)對(duì)象或傳遞給管道的參數(shù)这弧。改變這些狀態(tài)后， 你可以對(duì)命令進(jìn)行編碼以執(zhí)行管道虚汛。編碼器將所有狀態(tài)更改和命令參數(shù)寫(xiě)入命令緩沖區(qū)匾浪。
設(shè)置管道狀態(tài)和參數(shù)數(shù)據(jù)
設(shè)置要執(zhí)行命令的管道的管道狀態(tài)對(duì)象，然后為管道需要發(fā)送到add_arrays函數(shù)的任何參數(shù)設(shè)置數(shù)據(jù)卷哩。對(duì)于這個(gè)管道蛋辈，這意味著提供對(duì)三個(gè)緩沖區(qū)的引用。Metal自動(dòng)按照清單2中函數(shù)聲明中參數(shù)出現(xiàn)的順序?yàn)榫彌_區(qū)參數(shù)分配索引将谊，從0開(kāi)始冷溶。使用相同的索引提供參數(shù)。
[computeEncoder setComputePipelineState:_mAddFunctionPSO];
[computeEncoder setBuffer:_mBufferA offset:0 atIndex:0];
[computeEncoder setBuffer:_mBufferB offset:0 atIndex:1];
[computeEncoder setBuffer:_mBufferResult offset:0 atIndex:2];
還可以為每個(gè)參數(shù)指定偏移量尊浓。偏移量為0表示命令將從緩沖區(qū)開(kāi)始訪問(wèn)數(shù)據(jù)逞频。不管怎么樣，你可以使用一個(gè)緩沖區(qū)存儲(chǔ)多個(gè)參數(shù)眠砾，為每個(gè)參數(shù)指定偏移量虏劲。
你沒(méi)有為索引參數(shù)提供任何參數(shù)，因?yàn)閍dd_arrays函數(shù)定義它的值由GPU提供褒颈。
指定線程數(shù)和組織
接下來(lái)柒巫，決定要?jiǎng)?chuàng)建多少線程以及如何組織這些線程。Metal 可以創(chuàng)建1維谷丸，2維或3維網(wǎng)格堡掏。add_arrays函數(shù)使用一維數(shù)組，因此該示例創(chuàng)建了一個(gè)大小為（dataSize x 1 x 1）的一維網(wǎng)格刨疼，Metal從中生成0到dataSize-1之間的索引泉唁。
MTLSize gridSize = MTLSizeMake(arrayLength, 1, 1);
指定線程組大小
Metal 將網(wǎng)格細(xì)分為更小的網(wǎng)格，成為線程組揩慕。每個(gè)線程組單獨(dú)運(yùn)算亭畜。Metal 可以將線程組分配給GPU上的不同的處理元素以加快處理速度。您還需要確定為命令創(chuàng)建線程組的大小迎卤。
NSUInteger threadGroupSize = _mAddFunctionPSO.maxTotalThreadsPerThreadgroup;
if (threadGroupSize > arrayLength)
{
threadGroupSize = arrayLength;
}
MTLSize threadgroupSize = MTLSizeMake(threadGroupSize, 1, 1);
應(yīng)用程序向管道狀態(tài)對(duì)象請(qǐng)求最大可能的線程組拴鸵，如果該大小大于數(shù)據(jù)集的大小，則會(huì)收縮該線程組。 maxTotalThreadsPerThreadgrou屬性給一個(gè)線程組允許的最大的值劲藐，它根據(jù)用于創(chuàng)建管道狀態(tài)對(duì)象的函數(shù)的復(fù)雜性的不同而不同八堡。
編碼運(yùn)算命令以執(zhí)行線程
最后，對(duì)命令進(jìn)行編碼以調(diào)度線程網(wǎng)格聘芜。
[computeEncoder dispatchThreads:gridSize
threadsPerThreadgroup:threadgroupSize];
當(dāng)GPU執(zhí)行此命令時(shí)兄渺，它使用你先前設(shè)置的狀態(tài)和命令的參數(shù)來(lái)調(diào)度線程來(lái)執(zhí)行計(jì)算。你可以按照相同的步驟使用編碼器將多個(gè)計(jì)算命令編碼到計(jì)算過(guò)程中汰现，而無(wú)需執(zhí)行任何冗余步驟挂谍。例如，可以設(shè)置管道狀態(tài)對(duì)象一次服鹅，然后為要處理的每個(gè)緩沖區(qū)集合設(shè)置參數(shù)并編碼一個(gè)命令凳兵。
結(jié)束計(jì)算過(guò)程
當(dāng)沒(méi)有更多的命令添加到計(jì)算過(guò)程時(shí)，結(jié)束編碼過(guò)程以結(jié)束計(jì)算過(guò)程企软。
[computeEncoder endEncoding];
提交命令緩沖區(qū)以執(zhí)行其命令
通過(guò)將命令緩沖區(qū)提交到隊(duì)列來(lái)執(zhí)行命令緩沖區(qū)中的命令庐扫。
[commandBuffer commit];
命令隊(duì)列創(chuàng)建了命令緩沖區(qū)，因此提交緩沖區(qū)時(shí)總是將其放在該隊(duì)列上仗哨。提交命令緩沖區(qū)后形庭，Metal 異步的準(zhǔn)備要執(zhí)行的命令，然后調(diào)度命令緩沖區(qū)在GPU上執(zhí)行厌漂。在GPU執(zhí)行完命令緩沖區(qū)的所有命令后萨醒，Metal將命令緩沖區(qū)標(biāo)記為已完成。
等待運(yùn)算完成
當(dāng)GPU處理命令時(shí)苇倡，應(yīng)用可以做其他的工作富纸。這個(gè)樣本不需要做任何額外的工作，所以它只需要等待 命令緩沖區(qū)完成旨椒。
[commandBuffer waitUntilCompleted];
或者晓褪，為了在Metal 處理完所有命令時(shí)得到通知，向命令緩沖區(qū)添加一個(gè)完成處理程序(addCompletedHandler(_:))综慎，或者通過(guò)讀取命令緩沖區(qū)的狀態(tài)status屬性來(lái)檢查它的狀態(tài)涣仿。
從緩沖區(qū)讀取結(jié)果
命令緩沖區(qū)完成后，GPU的運(yùn)算結(jié)果被保存在輸出緩沖區(qū)中示惊，并且Metal 執(zhí)行任何必要的步驟以確保CPU可以看到他們好港。在真實(shí)的應(yīng)用中，可以從緩沖區(qū)中讀取結(jié)果并對(duì)其進(jìn)行處理米罚，例如在屏幕上顯示結(jié)果或?qū)⒔Y(jié)果寫(xiě)入文件钧汹。因?yàn)橛?jì)算只是用來(lái)說(shuō)明創(chuàng)建一個(gè)Metal 應(yīng)用程序的過(guò)程，該示例讀取存儲(chǔ)在輸出緩沖區(qū)中的值并進(jìn)行測(cè)試录择，以確保CPU和GPU計(jì)算的結(jié)果相同
- (void) verifyResults
{
float* a = _mBufferA.contents;
float* b = _mBufferB.contents;
float* result = _mBufferResult.contents;
for (unsigned long index = 0; index < arrayLength; index++)
{
if (result[index] != (a[index] + b[index]))
{
printf("Compute ERROR: index=%lu result=%g vs %g=a+b\n",
index, result[index], a[index] + b[index]);
assert(result[index] == (a[index] + b[index]));
}
}
printf("Compute results as expected\n");
}
