在GPU上執(zhí)行計(jì)算

使用Metal 查找GPU并對(duì)其進(jìn)行計(jì)算。

概述

在此示例中，您將學(xué)習(xí)所有Metal應(yīng)用程序中使用的基本任務(wù)，您將看到如何將用C編寫(xiě)的簡(jiǎn)單函數(shù)轉(zhuǎn)換為MetalShadingLanguage(MTL)，以便可以在GPU上運(yùn)行。您將找到一個(gè)GPU,通過(guò)創(chuàng)建管道來(lái)準(zhǔn)備要在其上運(yùn)行的MSL函數(shù)，并創(chuàng)建可供GPU訪問(wèn)的數(shù)據(jù)對(duì)象。要針對(duì)您的數(shù)據(jù)執(zhí)行管道，請(qǐng)創(chuàng)建命令緩沖區(qū)，將命令寫(xiě)入其中，然后將緩沖區(qū)提交到命令隊(duì)列中。Metal將命令發(fā)送到GPU進(jìn)行執(zhí)行。

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

為了說(shuō)明GPU編程，此應(yīng)用程序?qū)蓚€(gè)數(shù)組的相應(yīng)元素加在一起，然后將結(jié)果寫(xiě)入第三個(gè)數(shù)組。清單1顯示了一個(gè)用c編寫(xiě)的在CPU上執(zhí)行此計(jì)算的函數(shù)。它循環(huán)遍歷索引，每次循環(huán)迭代計(jì)算一個(gè)值。

清單1用C編寫(xiě)的數(shù)組加法

 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ì)算的，因此可以安全地同時(shí)計(jì)算這些值。要在GPU上執(zhí)行計(jì)算，您需要使用Metal Shading Language(MSL)重寫(xiě)此功能。MSL是c++的一種變體，專(zhuān)為GPU編程設(shè)計(jì)。在Metal中，在GPU上運(yùn)行的代碼稱(chēng)為著色器，因?yàn)闅v史上他們首先用于計(jì)算3D圖形中的顏色，清單2顯示了MSL中的著色器，該著色器使用執(zhí)行與清單1相同的計(jì)算，樣例項(xiàng)目在add.metal文件中定義了此函數(shù)。xcode會(huì)在應(yīng)用程序目標(biāo)中構(gòu)建所有.metal文件，并創(chuàng)建一個(gè)默認(rèn)的Metal庫(kù)，并將其嵌入到您的應(yīng)用程序中，您將在此示例后面看到如何加載默認(rèn)庫(kù)。

清單2用MSL編寫(xiě)的數(shù)組加法

 kernel void add_arrays( device const float *inA,
                          device const float *inB,
                        device const float *result,
              uint index [[thread_position_in_grid]]) {
  result[index] = inA[index] + inB[index];

清單1和清單2相似，但使用MSL版本有一些重要的區(qū)別。仔細(xì)查看清單2。
首先該函數(shù)添加kernel關(guān)鍵字，該關(guān)鍵字聲明該函數(shù)為：

一個(gè)公共GPU功能。公共功能是您應(yīng)用程序可以看到的唯一功能。其他著色器也不能調(diào)用公共功能呢。
一個(gè) compute 函數(shù)（也稱(chēng)為計(jì)算內(nèi)核），其執(zhí)行使用線程的網(wǎng)格中的并行計(jì)算。

請(qǐng)閱讀使用渲染管道渲染基元
以了解用于聲明公共圖形功能的其他功能關(guān)鍵字。

該函數(shù)使用device關(guān)鍵字聲明其三個(gè)參數(shù)，該參數(shù)表示這些指針位于地址空間中，MSL為內(nèi)存定義了幾個(gè)不相交的地址空間。每當(dāng)在MSL中聲明指針時(shí)，都必須提供device關(guān)鍵字來(lái)聲明其地址空間。使用地址空間聲明GPU可以讀取和寫(xiě)入的永久內(nèi)存。

清單2從清單1中刪除了for循環(huán)，因?yàn)樵摵瘮?shù)現(xiàn)在將被計(jì)算網(wǎng)格中的多個(gè)網(wǎng)格調(diào)用。此示例創(chuàng)建與數(shù)組尺寸完全匹配的一維線程網(wǎng)路，以便數(shù)組中的每個(gè)條目都有不同的線程計(jì)算。

為了替換以前由for循環(huán)提供的索引，該函數(shù)采用一個(gè)新的index參數(shù)，該參數(shù)帶有另一個(gè)MSL,thread_position_in_grid關(guān)鍵字,該關(guān)鍵字使用c++屬性語(yǔ)法指定,該關(guān)鍵字聲明Metal應(yīng)該為每個(gè)線程計(jì)算一個(gè)唯一索引，并在該參數(shù)中傳遞該索引。由于使用一維網(wǎng)格，因此將索引定義為標(biāo)量整數(shù)。即時(shí)刪除了循環(huán)，清單1和清單2扔使用同一行代碼將兩個(gè)數(shù)字加在一起。如果要將類(lèi)似的代碼從c或c++轉(zhuǎn)換為MSL,請(qǐng)以相同的方式用網(wǎng)格替換為循環(huán)邏輯.

尋找GPU

在您的應(yīng)用程序中，MTLDevice對(duì)象是GPU精簡(jiǎn)的抽象。您可以使用它與GPU進(jìn)行通信。Metal為每個(gè)GPU創(chuàng)建一個(gè)MTLDevice。您可以通過(guò)調(diào)用MTLCreateSystemDefaultDevice獲取默認(rèn)設(shè)備對(duì)象。在Mac上，Mac可以有多個(gè)GPU，在Metal ios 中，Metal 選擇其中一個(gè)GPU作為默認(rèn)GPU，然后返回該GPU的設(shè)備對(duì)象，在macos中，Metal提供了可用于所有設(shè)備對(duì)象的其他api，但是此示例僅適用默認(rèn)值。

Id<MTLDevice> device = MTLCreateSystemDefaultDevice();

初始化金屬對(duì)象

metal 將其他與GPU相關(guān)的實(shí)體（例如已編譯的著色器，內(nèi)存緩沖區(qū)或紋理）表示為對(duì)象。要?jiǎng)?chuàng)建這些特定GPU對(duì)象，請(qǐng)調(diào)用MTLDevice相關(guān)方法，或在MTLDevice上創(chuàng)建對(duì)象的方法。由設(shè)備對(duì)象直接或間接創(chuàng)建的所有對(duì)象只能與該設(shè)備對(duì)象一起使用。使用多個(gè)GPU的應(yīng)用程序?qū)⑹褂枚鄠€(gè)設(shè)備對(duì)象，并為每個(gè)對(duì)象創(chuàng)建相似的Metal對(duì)象層次結(jié)構(gòu)。

該示例應(yīng)用程序使用自定義類(lèi)MetalAddr來(lái)管理與GPU通信所需的對(duì)象。類(lèi)的初始化程序創(chuàng)建這些對(duì)象并將其存儲(chǔ)在其屬性中，該應(yīng)用程序創(chuàng)建此類(lèi)的實(shí)例，傳入Metal設(shè)備對(duì)象以用于創(chuàng)建輔助對(duì)象。該對(duì)象將對(duì)Metal對(duì)象保持強(qiáng)烈引用。直到完成執(zhí)行為止。

MetalAddr *adder = [[MetalAddr alloc】 init】；

在Metal 中，昂貴的初始化任務(wù)可以執(zhí)行一次，結(jié)果可以廉價(jià)地保留和使用，您很少需要在性能敏感的代碼中運(yùn)行此類(lèi)任務(wù)。

獲取金屬功能的參考

初始化程序要做的第一件事是加載函數(shù)并準(zhǔn)備使其在GPU上運(yùn)行。當(dāng)您構(gòu)建應(yīng)用程序時(shí)，xcode 會(huì) 編譯該函數(shù)并將其添加嵌入在應(yīng)用程序中的默認(rèn)Metal庫(kù)中，您可以使用MTLLibrary和MTLFuncation對(duì)象獲取有關(guān)Metal庫(kù)及其中包含的功能的信息。要獲取add_arrays代表該功能的MTLFuncation對(duì)象，請(qǐng)要求MTLDevice為默認(rèn)庫(kù)創(chuàng)建一個(gè)MTLLIbrary對(duì)象，然后向該庫(kù)要求一個(gè)代表著色器功能的MTLFuncation對(duì)象。

- (instancetype) initWithDevice:(id<MTLDevice>) device {

   self = [self init];
  if (self) {
        _mDevice = device;
    NSError *error = nil;
      id<MTLLibrary> defaultLibrary = [_mDevice newDefaultLibrary];
     if (defaultLibrary == nil ) {
        NSLog(@"Failed to find the default library");
        return nil;
    }
     id <MTLFuncation> addFunction = [defaultLibrary newFuncationWithName:@"add_arrays"];
        if (addFuncation == nil ) {
                 NSLog(@"Failed to find the adder Funcation");
                  return nil;
            }
   }
}

準(zhǔn)備金屬管道

函數(shù)對(duì)象是MSL函數(shù)的代理，但不是可執(zhí)行代碼。您可以通過(guò)創(chuàng)建管道將函數(shù)轉(zhuǎn)換為可執(zhí)行代碼。管道指定GPU執(zhí)行以完成特定任務(wù)的步驟。在Metal中，管道由管道狀態(tài)對(duì)象表示。由于此示例使用計(jì)算功能，因此該應(yīng)用程序?qū)?chuàng)建一個(gè)MTLComputePipelineState對(duì)象。

_mAddFuncationPSO = [_mDevice newComputePipelineStateWithFuncation:addFuncation error:&error];

計(jì)算管道運(yùn)行單個(gè)計(jì)算功能，可以選擇在運(yùn)行功能之前處理輸入數(shù)據(jù),然后再運(yùn)行輸出數(shù)據(jù)。

創(chuàng)建管道狀態(tài)對(duì)象時(shí)，設(shè)備對(duì)象將完成針對(duì)該特定GPU的功能編譯。此示例同步創(chuàng)建管道狀態(tài)對(duì)象，并將其直接返回給應(yīng)用程序。由于編輯確實(shí)需要一段時(shí)間。因此請(qǐng)避免在對(duì)性能敏感的代碼中同步創(chuàng)建管道狀態(tài)對(duì)象。

筆記
到目前為止,您在代碼中看到的Metal返回的所有對(duì)象都將作為符合協(xié)議的對(duì)象返回。Metal使用協(xié)議定義了大多數(shù)特定于GPU的對(duì)象，以抽象出底層的實(shí)現(xiàn)類(lèi)。這對(duì)于不同GPU可能有所不同。Metal使用類(lèi)定義了獨(dú)立于GPU的對(duì)象。任何給定的Metal協(xié)議的參考文檔都明確說(shuō)明了您是否可以在您是否可以在您的應(yīng)用中實(shí)現(xiàn)該協(xié)議。

創(chuàng)建命令隊(duì)列

要將工作發(fā)送到GPU,您需要一個(gè)命令隊(duì)列。 Metal 使用命令隊(duì)列來(lái)調(diào)度命令。 通過(guò)詢(xún)問(wèn)一個(gè)MTLDevice來(lái)創(chuàng)建一個(gè)命令隊(duì)列.

_mCommandQueue = [_mDevice newCommandQueue];

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

初始化基本的metal對(duì)象后,您將加載數(shù)據(jù)以供GPU執(zhí)行。此任務(wù)對(duì)性能的要求不高，但在應(yīng)用程序啟動(dòng)初期仍然有用。

GPU 可以擁有自己的專(zhuān)用內(nèi)存,也可以與操作系統(tǒng)共享內(nèi)存.

Metal 和操作系統(tǒng)內(nèi)核需要執(zhí)行其它工作，才能讓您將數(shù)據(jù)存儲(chǔ)在內(nèi)存并使這些數(shù)據(jù)可用于GPU。Metal 使用資源MTLResource對(duì)象抽象了此內(nèi)存管理。資源是運(yùn)行命令時(shí)GPU可以訪問(wèn)的內(nèi)存分配。使用MTLDevice為它的GPU創(chuàng)建資源。

該示例應(yīng)用程序?qū)?chuàng)建三個(gè)緩沖區(qū)，并使用隨機(jī)數(shù)據(jù)填充前三個(gè)緩沖區(qū)。第三個(gè)緩沖區(qū)是存儲(chǔ)結(jié)果的位置。

 _mBufferA = [_mDevice newBufferWIthLength:bufferSize options:MTLResourceStorageShared];
_mBufferB = [_mDevice newBufferWithLength:bufferSize options:MTLResourceStorageShared];
_mBufferResult = [_mDevice newBufferWithLength:bufferSize options:MTLResourceStoreShared];

[self generateRandomFloatData：_mBufferA];
[self generateRandomFloatData:  _mBufferB];

此示例中的資源是MTLBuffer對(duì)象，他們是沒(méi)有預(yù)定義格式的內(nèi)存分配。Metal將每個(gè)緩沖區(qū)作為不透明的字節(jié)集合進(jìn)行管理。但是，在著色器中使用緩沖區(qū)時(shí)，需要指定格式。這意味著您的著色器和應(yīng)用程序需要就來(lái)回傳遞的任何數(shù)據(jù)的格式達(dá)成共識(shí)。

分配緩沖區(qū)時(shí)，您將提供一種存儲(chǔ)模式，以確定某些性能特征以及CPU或GPU是可以訪問(wèn)它。該示例應(yīng)用程序使用GPU和CPU都可以訪問(wèn)的共享內(nèi)存 StorageModeShared。

為了用隨機(jī)數(shù)據(jù)填充緩沖區(qū)，應(yīng)用程序獲取指向緩沖區(qū)內(nèi)存的指針，并將數(shù)據(jù)寫(xiě)入CPU。清單2中的函數(shù)將其參數(shù)聲明為浮點(diǎn)數(shù)數(shù)組，因此您提供了相同格式的緩沖區(qū)：add_arrays


 - (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)建命令緩沖區(qū)

要求命令隊(duì)列創(chuàng)建命令緩沖區(qū)

 id <MTLCommandbuffer> commandbuffer =  [_commandQueue commandBuffer]；

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

要將命令寫(xiě)入命令緩沖區(qū)，可以對(duì)要編碼的特定類(lèi)型的命令使用命令編碼器，該示例創(chuàng)建了一個(gè)計(jì)算命令編碼器，該編碼器對(duì)一個(gè)計(jì)算通道進(jìn)行編碼。計(jì)算過(guò)程包含執(zhí)行計(jì)算管道的命令列表，每個(gè)計(jì)算命令使GPU創(chuàng)建線程網(wǎng)格以在GPU上執(zhí)行

id<MTLComputeCommandEncoder> commandEncoder = [commandBuffer makeComputeEncoder];

要對(duì)命令進(jìn)行編碼，您需要在編碼器上進(jìn)行一系列方法調(diào)用。一些方法設(shè)置狀態(tài)信息，例如管道狀態(tài)對(duì)象（pso）或要傳遞給管道的參數(shù)。進(jìn)行狀態(tài)變更后，對(duì)命令進(jìn)行編碼以執(zhí)行管道。編碼器將所有狀態(tài)更改和命令參數(shù)寫(xiě)入命令緩沖區(qū)。

9e65eb88-9081-40fc-b2e4-4c72877bfcd4.png

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

設(shè)置您要命令執(zhí)行的管道的管道狀態(tài)對(duì)象。然后為管道需要發(fā)送到函數(shù)中的所有參數(shù)設(shè)置數(shù)據(jù)。對(duì)于此管道，這意味著提供三個(gè)緩沖區(qū)的引用。Metal 會(huì)按照參數(shù)在清單2的add_arrays函數(shù)聲明中出現(xiàn)的順序自動(dòng)為緩沖區(qū)參數(shù)分配索引。您使用相同的索引提供參數(shù).

[computeEncoder setcomputeState:_mAaddFunctionPSO];
[computeEncoder setBuffer:_mBufferA offset:0 atIndex:0];
[computeEncoder setButter:_mBufferB offset:0 atIndex:1];
[computeEncoder setBuffer:_result offset:0 atIndex:2];

您還可以為每個(gè)參數(shù)指定一個(gè)偏移量。偏移為0 表示命令將從緩沖區(qū)的開(kāi)頭訪問(wèn)數(shù)據(jù)。但是，您可以使用一個(gè)緩沖區(qū)來(lái)存儲(chǔ)多個(gè)參數(shù)，并為每個(gè)參數(shù)指定一個(gè)偏移量。

您無(wú)需為index 參數(shù)指定任何數(shù)據(jù)，因?yàn)樵揳dd_arrays函數(shù)將其值定義由GPU提供。

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

接下來(lái)，確定要?jiǎng)?chuàng)建多少個(gè)線程以及如何組織這些線程。金屬可以創(chuàng)建1D、2D、3D網(wǎng)格。該函數(shù)使用一維數(shù)組，因此示例將創(chuàng)建一個(gè)大小為（array_length x 1x 1)的一維網(wǎng)格，Metal將根據(jù)該網(wǎng)格生成介于0 到 array_length -1 之間的索引。

 MTLSize gridSize = MTLSize(array_length,1,1)；

指定線程組大小

金屬將網(wǎng)格分為更小的網(wǎng)格，稱(chēng)為線程組。每個(gè)線程組是單獨(dú)計(jì)算的，metal可以將線程組調(diào)度到GPU上的不同處理元素，以加快處理速度。您還需要確定為命令創(chuàng)建線程組的大小。

NSUinteger threadGroupSize = _mAddFunctionPSO.maxTotalThreadsPerThreadGroup；
if （threadGroupSize > arrayLength） {
   threadGroupSize = arrayLength;
}
MTLSize threadgroupsize = MTLSize(threadGroupSize,1,1);

該應(yīng)用程序向管道狀態(tài)對(duì)象詢(xún)問(wèn)最大可能的線程組。如果該大小大于數(shù)據(jù)集的大小，則將其壓縮，該maxTotalThreadsPerThreadGroup屬性提供線程組中允許的最大線程數(shù)，該數(shù)量根據(jù)用于創(chuàng)建管道狀態(tài)對(duì)象的函數(shù)的復(fù)雜性而變化。

編碼compute命令以執(zhí)行線程

最后，對(duì)命令進(jìn)行編碼以調(diào)度線程網(wǎng)格。

 [computeEncoder dispathThreads:gridSize threadsPerTheadGroup:threadgroupSize];

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

您可以使用編碼器按照相同的步驟將多個(gè)計(jì)算命令編碼到計(jì)算過(guò)程中，而無(wú)需執(zhí)行任何冗余步驟，例如，您可以設(shè)置一次管道狀態(tài)的對(duì)象，然后為每個(gè)要處理的緩沖區(qū)集合設(shè)置參數(shù)并編碼一個(gè)命令。

結(jié)束計(jì)算階段

如果沒(méi)有更多命令要添加到計(jì)算過(guò)程中，則結(jié)束編碼過(guò)程以結(jié)束計(jì)算過(guò)程。

[computerEncoder encoding];

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

通過(guò)將命令緩沖區(qū)提交到運(yùn)行命令緩沖區(qū)中的命令。

[commBuffer commit ];

命令隊(duì)列創(chuàng)建了緩沖區(qū)，因此提交緩沖區(qū)總是將其放置在該隊(duì)列中。
GPU執(zhí)行命令緩沖區(qū)中的所有命令后，Metal將命令緩沖區(qū)標(biāo)記為已完成。

等待計(jì)算完成

在GPU處理命令時(shí)，您的應(yīng)用程序可以執(zhí)行其他工作。該示例不需要執(zhí)行任何其他工作。因此只需等待命令緩沖區(qū)完成即可。

[computeBuffer waitUntilCompete];

或者要在metal處理完成所有命令時(shí)得到通知，請(qǐng)將完成處理程序addComputeHander(_:)添加到命令緩沖區(qū)，或者通過(guò)讀取status屬性來(lái)檢查命令緩沖區(qū)的狀態(tài)。

從緩沖區(qū)讀取結(jié)果

命令緩沖區(qū)完成后，GPU的計(jì)算將存儲(chǔ)在輸出緩沖區(qū)中，metal將執(zhí)行任務(wù)必要的步驟以確保cpu可以看到他們。在真實(shí)的應(yīng)用程序中，您需要從緩沖區(qū)讀取結(jié)果并對(duì)結(jié)果進(jìn)行處理。

色偷偷精品伊人,欧洲久久精品,欧美综合婷婷骚逼,国产AV主播,国产最新探花在线,九色在线视频一区,伊人大交九欧美,1769亚洲,黄色成人av

Metal 基本任務(wù)和概念 - 01

Metal 基本任務(wù)和概念 - 01

在GPU上執(zhí)行計(jì)算

概述

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

尋找GPU

初始化金屬對(duì)象

獲取金屬功能的參考

準(zhǔn)備金屬管道

創(chuàng)建命令隊(duì)列

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

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

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

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

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

指定線程組大小

編碼compute命令以執(zhí)行線程

結(jié)束計(jì)算階段

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

等待計(jì)算完成

從緩沖區(qū)讀取結(jié)果

相關(guān)閱讀更多精彩內(nèi)容

友情鏈接更多精彩內(nèi)容

色偷偷精品伊人,欧洲久久精品,欧美综合婷婷骚逼,国产AV主播,国产最新探花在线,九色在线视频一区,伊人大交九 欧美,1769亚洲,黄色成人av

Metal 基本任務(wù)和概念 - 01

在GPU上執(zhí)行計(jì)算

概述

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

尋找GPU

初始化金屬對(duì)象

獲取金屬功能的參考

準(zhǔn)備金屬管道

創(chuàng)建命令隊(duì)列

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

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

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

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

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

指定線程組大小

編碼compute命令以執(zhí)行線程

結(jié)束計(jì)算階段

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

等待計(jì)算完成

從緩沖區(qū)讀取結(jié)果

相關(guān)閱讀更多精彩內(nèi)容

友情鏈接更多精彩內(nèi)容

色偷偷精品伊人,欧洲久久精品,欧美综合婷婷骚逼,国产AV主播,国产最新探花在线,九色在线视频一区,伊人大交九欧美,1769亚洲,黄色成人av