在GPU上执行计算
使用Metal 查找GPU并对其进行计算。
概述
在此示例中，您将学习所有Metal应用程序中使用的基本任务，您将看到如何将用C编写的简单函数转换为MetalShadingLanguage(MTL)，以便可以在GPU上运行。您将找到一个GPU,通过创建管道来准备要在其上运行的MSL函数，并创建可供GPU访问的数据对象。要针对您的数据执行管道，请创建命令缓冲区，将命令写入其中，然后将缓冲区提交到命令队列中。Metal将命令发送到GPU进行执行。
编写GPU函数以执行计算
为了说明GPU编程，此应用程序将两个数组的相应元素加在一起，然后将结果写入第三个数组。清单1显示了一个用c编写的在CPU上执行此计算的函数。它循环遍历索引，每次循环迭代计算一个值。
清单1用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上执行计算，您需要使用Metal Shading Language(MSL)重写此功能。MSL是c++的一种变体，专为GPU编程设计。在Metal中，在GPU上运行的代码称为着色器，因为历史上他们首先用于计算3D图形中的颜色，清单2显示了MSL中的着色器，该着色器使用执行与清单1相同的计算，样例项目在add.metal文件中定义了此函数。xcode会在应用程序目标中构建所有.metal文件，并创建一个默认的Metal库，并将其嵌入到您的应用程序中，您将在此示例后面看到如何加载默认库。
清单2用MSL编写的数组加法
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版本有一些重要的区别。仔细查看清单2。
首先该函数添加kernel关键字，该关键字声明该函数为：
一个公共GPU功能。公共功能是您应用程序可以看到的唯一功能。其他着色器也不能调用公共功能呢。
一个 compute 函数（也称为计算内核），其执行使用线程的网格中的并行计算。
请阅读使用渲染管道渲染基元
以了解用于声明公共图形功能的其他功能关键字。
该函数使用device关键字声明其三个参数，该参数表示这些指针位于地址空间中，MSL为内存定义了几个不相交的地址空间。每当在MSL中声明指针时，都必须提供device关键字来声明其地址空间。使用地址空间声明GPU可以读取和写入的永久内存。
清单2从清单1中删除了for循环，因为该函数现在将被计算网格中的多个网格调用。此示例创建与数组尺寸完全匹配的一维线程网路，以便数组中的每个条目都有不同的线程计算。
为了替换以前由for循环提供的索引，该函数采用一个新的index参数，该参数带有另一个MSL,thread_position_in_grid关键字,该关键字使用c++属性语法指定,该关键字声明Metal应该为每个线程计算一个唯一索引，并在该参数中传递该索引。由于使用一维网格，因此将索引定义为标量整数。即时删除了循环，清单1和清单2扔使用同一行代码将两个数字加在一起。如果要将类似的代码从c或c++转换为MSL,请以相同的方式用网格替换为循环逻辑.
寻找GPU
在您的应用程序中，MTLDevice对象是GPU精简的抽象。您可以使用它与GPU进行通信。Metal为每个GPU创建一个MTLDevice。您可以通过调用MTLCreateSystemDefaultDevice获取默认设备对象。在Mac上，Mac可以有多个GPU，在Metal ios 中，Metal 选择其中一个GPU作为默认GPU，然后返回该GPU的设备对象，在macos中，Metal提供了可用于所有设备对象的其他api，但是此示例仅适用默认值。
Id<MTLDevice> device = MTLCreateSystemDefaultDevice();
初始化金属对象
metal 将其他与GPU相关的实体（例如已编译的着色器，内存缓冲区或纹理）表示为对象。要创建这些特定GPU对象，请调用MTLDevice相关方法，或在MTLDevice上创建对象的方法。由设备对象直接或间接创建的所有对象只能与该设备对象一起使用。使用多个GPU的应用程序将使用多个设备对象，并为每个对象创建相似的Metal对象层次结构。
该示例应用程序使用自定义类MetalAddr来管理与GPU通信所需的对象。类的初始化程序创建这些对象并将其存储在其属性中，该应用程序创建此类的实例，传入Metal设备对象以用于创建辅助对象。该对象将对Metal对象保持强烈引用。直到完成执行为止。
MetalAddr *adder = [[MetalAddr alloc】 init】；
在Metal 中，昂贵的初始化任务可以执行一次，结果可以廉价地保留和使用，您很少需要在性能敏感的代码中运行此类任务。
获取金属功能的参考
初始化程序要做的第一件事是加载函数并准备使其在GPU上运行。当您构建应用程序时，xcode 会 编译该函数并将其添加嵌入在应用程序中的默认Metal库中，您可以使用MTLLibrary和MTLFuncation对象获取有关Metal库及其中包含的功能的信息。要获取add_arrays代表该功能的MTLFuncation对象，请要求MTLDevice为默认库创建一个MTLLIbrary对象，然后向该库要求一个代表着色器功能的MTLFuncation对象。
- (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;
}
}
}
准备金属管道
函数对象是MSL函数的代理，但不是可执行代码。您可以通过创建管道将函数转换为可执行代码。管道指定GPU执行以完成特定任务的步骤。 在Metal中，管道由管道状态对象表示。由于此示例使用计算功能，因此该应用程序将创建一个MTLComputePipelineState对象。
_mAddFuncationPSO = [_mDevice newComputePipelineStateWithFuncation:addFuncation error:&error];
计算管道运行单个计算功能，可以选择在运行功能之前处理输入数据,然后再运行输出数据。
创建管道状态对象时，设备对象将完成针对该特定GPU的功能编译。此示例同步创建管道状态对象，并将其直接返回给应用程序。由于编辑确实需要一段时间。因此请避免在对性能敏感的代码中同步创建管道状态对象。
笔记
到目前为止,您在代码中看到的Metal返回的所有对象都将作为符合协议的对象返回。Metal使用协议定义了大多数特定于GPU的对象，以抽象出底层的实现类。这对于不同GPU可能有所不同。Metal使用类定义了独立于GPU的对象。任何给定的Metal协议的参考文档都明确说明了您是否可以在您是否可以在您的应用中实现该协议。
创建命令队列
要将工作发送到GPU,您需要一个命令队列。 Metal 使用命令队列来调度命令。 通过询问一个MTLDevice来创建一个命令队列.
_mCommandQueue = [_mDevice newCommandQueue];
创建数据缓冲区并加载数据
初始化基本的metal对象后,您将加载数据以供GPU执行。此任务对性能的要求不高，但在应用程序启动初期仍然有用。
GPU 可以拥有自己的专用内存,也可以与操作系统共享内存.
Metal 和操作系统内核需要执行其它工作，才能让您将数据存储在内存并使这些数据可用于GPU。Metal 使用资源MTLResource对象抽象了此内存管理。资源是运行命令时GPU可以访问的内存分配。使用MTLDevice为它的GPU创建资源。
该示例应用程序将创建三个缓冲区，并使用随机数据填充前三个缓冲区。第三个缓冲区是存储结果的位置。
_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对象，他们是没有预定义格式的内存分配。Metal将每个缓冲区作为不透明的字节集合进行管理。但是，在着色器中使用缓冲区时，需要指定格式。这意味着您的着色器和应用程序需要就来回传递的任何数据的格式达成共识。
分配缓冲区时，您将提供一种存储模式，以确定某些性能特征以及CPU或GPU是可以访问它。该示例应用程序使用GPU和CPU都可以访问的共享内存 StorageModeShared。
为了用随机数据填充缓冲区，应用程序获取指向缓冲区内存的指针，并将数据写入CPU。清单2中的函数将其参数声明为浮点数数组，因此您提供了相同格式的缓冲区：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);
}
}
创建命令缓冲区
要求命令队列创建命令缓冲区
id <MTLCommandbuffer> commandbuffer = [_commandQueue commandBuffer]；
创建命令编码器
要将命令写入命令缓冲区，可以对要编码的特定类型的命令使用命令编码器，该示例创建了一个计算命令编码器，该编码器对一个计算通道进行编码。计算过程包含执行计算管道的命令列表，每个计算命令使GPU创建线程网格以在GPU上执行
id<MTLComputeCommandEncoder> commandEncoder = [commandBuffer makeComputeEncoder];
要对命令进行编码，您需要在编码器上进行一系列方法调用。一些方法设置状态信息，例如管道状态对象（pso）或要传递给管道的参数。进行状态变更后，对命令进行编码以执行管道。编码器将所有状态更改和命令参数写入命令缓冲区。
设置管道状态和参数数据
设置您要命令执行的管道的管道状态对象。然后为管道需要发送到函数中的所有参数设置数据。对于此管道，这意味着提供三个缓冲区的引用。Metal 会按照参数在清单2的add_arrays函数声明中出现的顺序自动为缓冲区参数分配索引。您使用相同的索引提供参数.
[computeEncoder setcomputeState:_mAaddFunctionPSO];
[computeEncoder setBuffer:_mBufferA offset:0 atIndex:0];
[computeEncoder setButter:_mBufferB offset:0 atIndex:1];
[computeEncoder setBuffer:_result offset:0 atIndex:2];
您还可以为每个参数指定一个偏移量。偏移为0 表示命令将从缓冲区的开头访问数据。但是，您可以使用一个缓冲区来存储多个参数，并为每个参数指定一个偏移量。
您无需为index 参数指定任何数据，因为该add_arrays函数将其值定义由GPU提供。
指定线程数和组织
接下来，确定要创建多少个线程以及如何组织这些线程。金属可以创建1D、2D、3D网格。该函数使用一维数组，因此示例将创建一个大小为（array_length x 1x 1)的一维网格，Metal将根据该网格生成介于0 到 array_length -1 之间的索引。
MTLSize gridSize = MTLSize(array_length,1,1)；
指定线程组大小
金属将网格分为更小的网格，称为线程组。每个线程组是单独计算的，metal可以将线程组调度到GPU上的不同处理元素，以加快处理速度。您还需要确定为命令创建线程组的大小。
NSUinteger threadGroupSize = _mAddFunctionPSO.maxTotalThreadsPerThreadGroup；
if （threadGroupSize > arrayLength） {
threadGroupSize = arrayLength;
}
MTLSize threadgroupsize = MTLSize(threadGroupSize,1,1);
该应用程序向管道状态对象询问最大可能的线程组。如果该大小大于数据集的大小，则将其压缩，该maxTotalThreadsPerThreadGroup属性提供线程组中允许的最大线程数，该数量根据用于创建管道状态对象的函数的复杂性而变化。
编码compute命令以执行线程
最后，对命令进行编码以调度线程网格。
[computeEncoder dispathThreads:gridSize threadsPerTheadGroup:threadgroupSize];
GPU执行命令时，它将使用您先前设置的状态和命令的参数来分派线程以执行计算。
您可以使用编码器按照相同的步骤将多个计算命令编码到计算过程中，而无需执行任何冗余步骤，例如，您可以设置一次管道状态的对象，然后为每个要处理的缓冲区集合设置参数并编码一个命令。
结束计算阶段
如果没有更多命令要添加到计算过程中，则结束编码过程以结束计算过程。
[computerEncoder encoding];
提交命令缓冲区以执行命令
通过将命令缓冲区提交到运行命令缓冲区中的命令。
[commBuffer commit ];
- 命令队列创建了缓冲区，因此提交缓冲区总是将其放置在该队列中。
- GPU执行命令缓冲区中的所有命令后，Metal将命令缓冲区标记为已完成。
等待计算完成
在GPU处理命令时，您的应用程序可以执行其他工作。该示例不需要执行任何其他工作。因此只需等待命令缓冲区完成即可。
[computeBuffer waitUntilCompete];
或者要在metal处理完成所有命令时得到通知，请将完成处理程序addComputeHander(_:)添加到命令缓冲区，或者通过读取status属性来检查命令缓冲区的状态。
从缓冲区读取结果
命令缓冲区完成后，GPU的计算将存储在输出缓冲区中，metal将执行任务必要的步骤以确保cpu可以看到他们。在真实的应用程序中，您需要从缓冲区读取结果并对结果进行处理。
