本文是在苹果文档中学习总结的
使用Metal查找GPU并对其执行计算。
将介绍所有Metal应用程序中使用的基本任务：
将用C编写的简单函数转换为金属着色语言（MSL），以便它可以在GPU上运行
1、寻找GPU
2、通过创建管道准备MSL函数以在GPU上运行
3、创建GPU可访问的内存分配以保存数据
4、创建命令缓冲区并编码GPU命令以操纵数据
5、将缓冲区提交到命令队列以使GPU执行编码命令
编写GPU函数以执行计算
为了说明GPU编程，这个应用程序将两个数组的相应元素添加到一起，将结果写入第三个数组。清单1显示了一个在CPU上执行此计算的函数，用C语言编写。它循环遍历索引，计算循环的每次迭代一个值。
清单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];
}
}
const ****float *** p是指p是常量指针，指向float常量或者变量，但是可以改变指向；float *** **const **p 代表p可以指向变量，但是声明指定后就不能再改变
每个值都是独立计算的，因此可以同时安全地计算这些值。要在GPU上执行计算，需要在金属着色语言（MSL）中重写此功能。MSL是为GPU编程而设计的C ++的变体。在Metal中，在GPU上运行的代码称为着色器，因为历史上它们首先用于计算3D图形中的颜色。清单2显示了MSL中的着色器，它执行与清单1相同的计算。项目在文件中定义了此函数。Xcode构建应用程序目标中的所有文件，并创建一个默认的Metal库，它嵌入到您的应用程序中。将在本示例的后面看到如何加载默认库。add.metal.metal
清单2以MSL编写的数组添加
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类似，但MSL版本有一些重要的区别。仔细看看清单2。
首先，该函数添加kernel关键字，声明该函数是：
一个公共GPU功能。公共函数是您的应用可以看到的唯一功能。其他着色器函数也无法调用公共函数。
甲compute函数（也称为计算内核），其执行使用线程的网格中的并行计算。
请参阅使用渲染管道渲染基元以了解用于声明公共图形函数的其他函数关键字。
该函数使用关键字声明其三个参数，该关键字表示这些指针位于地址空间中。MSL为内存定义了几个不相交的地址空间。每当在MSL中声明指针时，都必须提供一个关键字来声明其地址空间。使用地址空间声明GPU可以读取和写入的持久内存。add_arraysdevicedevicedevice
清单2从清单1中删除了for循环，因为该函数现在将由计算网格中的多个线程调用。此示例创建一个完全匹配数组维度的一维线程网格，以便数组中的每个条目由不同的线程计算。
要替换先前由for循环提供的索引，该函数将index使用另一个MSL关键字获取一个新参数，该参数使用C ++属性语法指定。此关键字声明Metal应为每个线程计算唯一索引，并在此参数中传递该索引。因为使用1D网格，索引被定义为标量整数。即使删除了循环，清单1和清单2也使用相同的代码行将两个数字相加。如果要将类似的代码从C或C ++转换为MSL，请以相同的方式用网格替换循环逻辑。thread_position_in_gridadd_arrays
找一个GPU
在您的应用中，MTLDevice对象是GPU的精简抽象; 你用它来与GPU通信。Metal MTLDevice为每个GPU 创建一个。您可以通过调用获取默认设备对象 。在macOS中，Mac可以有多个GPU，Metal选择其中一个GPU作为默认值并返回该GPU的设备对象。在macOS中，Metal提供了可用于检索所有设备对象的其他API，但此示例仅使用默认值。MTLCreateSystemDefaultDevice
id<MTLDevice> device = MTLCreateSystemDefaultDevice();
初始化metal物体
Metal将其他与GPU相关的实体（如编译着色器，内存缓冲区和纹理）表示为对象。要创建这些特定于GPU的对象，可以在a上调用方法，MTLDevice或者在由a创建的对象上调用方法MTLDevice。由设备对象直接或间接创建的所有对象仅可用于该设备对象。使用多个GPU的应用程序将使用多个设备对象，并为每个设备创建类似的Metal对象层次结构。
示例应用程序使用自定义类来管理与GPU通信所需的对象。类的初始化程序创建这些对象并将它们存储在其属性中。该应用程序创建此类的实例，传入Metal设备对象以用于创建辅助对象。该对象保持对Metal对象的强引用，直到它完成执行。MetalAdderMetalAdder
MetalAdder* adder = [[MetalAdder alloc] initWithDevice:device];
在Metal中，昂贵的初始化任务可以运行一次，结果保留并且使用成本低廉。您很少需要在性能敏感的代码中运行此类任务。
获取金属功能的参考
初始化程序所做的第一件事是加载函数并准备它在GPU上运行。在构建应用程序时，Xcode会编译该函数并将其添加到它嵌入应用程序的默认Metal库中。您可以使用和对象获取有关Metal库及其中包含的函数的信息。要获取表示函数的对象，请要求为默认库创建对象，然后向库请求表示着色器函数的对象。add_arraysMTLLibraryMTLFunctionadd
- (instancetype) initWithDevice: (id<MTLDevice>) device
{
self = [super init];
if (self)
{
_mDevice = device;
NSError* error = nil;
//在项目中加载具有.metal文件扩展名的明暗器文件
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;
}
// 创建计算管道状态对象。
_mAddFunctionPSO = [_mDevice newComputePipelineStateWithFunction: addFunction error:&error];
if (_mAddFunctionPSO == nil)
{
// 如果启用了金属API验证，您可以了解有关about what
//出了问题。（运行调试生成时，默认情况下启用金属API验证is run
// 来自xcode）
NSLog(@"Failed to created pipeline state object, error %@.", error);
return nil;
}
_mCommandQueue = [_mDevice newCommandQueue];
if (_mCommandQueue == nil)
{
NSLog(@"Failed to find the command queue.");
return nil;
}
}
return self;
}
准备金属管道
函数对象是MSL函数的代理，但它不是可执行代码。您可以通过创建管道将该函数转换为可执行代码。管道指定GPU执行以完成特定任务的步骤。在Metal中，管道由管道状态对象表示。由于此示例使用计算功能，因此应用程序会创建一个对象。MTLComputePipelineState
_mAddFunctionPSO = [_mDevice newComputePipelineStateWithFunction: addFunction error:&error];
计算管道运行单个计算功能，可选地在运行函数之前操纵输入数据，然后运行输出数据。
创建管道状态对象时，设备对象将完成为此特定GPU编译的功能。此示例同步创建管道状态对象并将其直接返回给应用程序。因为编译确实需要一段时间，所以避免在性能敏感的代码中同步创建管道状态对象。
注意
Metal到目前为止您看到的代码中返回的所有对象都将作为符合协议的对象返回。Metal使用协议来定义大多数特定于GPU的对象，以抽象掉底层实现类，这些类可能因不同的GPU而异。Metal使用类定义与GPU无关的对象。任何给定的Metal协议的参考文档都清楚地表明您是否可以在应用程序中实现该协议。
创建一个命令队列
要将工作发送到GPU，您需要一个命令队列。Metal使用命令队列来安排命令。通过询问一个MTLDevice来创建一个命令队列。
_mCommandQueue = [_mDevice newCommandQueue];
创建数据缓冲区和加载数据
初始化基本Metal对象后，加载要执行的GPU数据。此任务的性能不太重要，但在应用程序启动的早期仍然有用。
GPU可以拥有自己的专用内存，也可以与操作系统共享内存。Metal和操作系统内核需要执行额外的工作，以便将数据存储在内存中并使数据可供GPU使用。Metal使用资源对象抽象此内存管理。（MTLResource）。资源是GPU在运行命令时可以访问的内存分配。使用a MTLDevice为其GPU创建资源。
示例应用程序创建三个缓冲区，并使用随机数据填充前两个缓冲区。第三个缓冲区将存储其结果。add_arrays
- (void) prepareData
{
// 分配三个缓冲区来保存初始数据和结果。
_mBufferA = [_mDevice newBufferWithLength:bufferSize options:MTLResourceStorageModeShared];
_mBufferB = [_mDevice newBufferWithLength:bufferSize options:MTLResourceStorageModeShared];
_mBufferResult = [_mDevice newBufferWithLength:bufferSize options:MTLResourceStorageModeShared];
[self generateRandomFloatData:_mBufferA];
[self generateRandomFloatData:_mBufferB];
}
此示例中的资源是（MTLBuffer）对象，它们是没有预定义格式的内存分配。Metal将每个缓冲区管理为不透明的字节集合。但是，在着色器中使用缓冲区时指定格式。这意味着您的着色器和您的应用需要就来回传递的任何数据的格式达成一致。
分配缓冲区时，您需要提供存储模式以确定其某些性能特征以及CPU或GPU是否可以访问它。示例应用程序使用共享内存（），CPU和GPU都可以访问它。MTLResourceStorageMode
要使用随机数据填充缓冲区，应用程序将获取指向缓冲区内存的指针，并在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 = [_mCommandQueue commandBuffer];
创建一个命令编码器
要将命令写入命令缓冲区，可以使用命令编码器来处理要编码的特定命令类型。此示例创建一个计算命令编码器，用于编码计算传递。计算传递包含执行计算管道的命令列表。每个计算命令都会使GPU创建一个在GPU上执行的线程网格。
id<MTLComputeCommandEncoder> computeEncoder = [commandBuffer computeCommandEncoder];
设置管道状态和参数数据
设置要执行命令的管道的管道状态对象。然后为管道需要发送到函数的任何参数设置数据。对于此管道，这意味着提供对三个缓冲区的引用。Metal以参数出现在清单2中的函数声明中的顺序自动为缓冲区参数指定索引，从。开始。您使用相同的索引提供参数。add_arrays0
[computeEncoder setComputePipelineState:_mAddFunctionPSO];
[computeEncoder setBuffer:_mBufferA offset:0 atIndex:0];
[computeEncoder setBuffer:_mBufferB offset:0 atIndex:1];
[computeEncoder setBuffer:_mBufferResult offset:0 atIndex:2];
您还可以为每个参数指定偏移量。偏移0意味着命令将从缓冲区的开头访问数据。但是，您可以使用一个缓冲区来存储多个参数，为每个参数指定偏移量。
您没有为index参数指定任何数据，因为该函数将其值定义为由GPU提供。add_arrays
指定线程计数和组织
接下来，确定要创建的线程数以及如何组织这些线程。Metal可以创建1D，2D或3D网格。该函数使用一维数组，因此样本创建一个大小为1 x的网格（x 1 x 1），Metal从该网格生成0到-1 之间的索引。add_arraysdataSizedataSize
MTLSize gridSize = MTLSizeMake(arrayLength, 1, 1);
指定线程组大小
Metal将网格细分为称为线程组的较小网格。每个线程组都是单独计算的。Metal可以将线程组分派到GPU上的不同处理元素，以加快处理速度。您还需要确定为命令创建线程组的大小。
NSUInteger threadGroupSize = _mAddFunctionPSO.maxTotalThreadsPerThreadgroup;
if (threadGroupSize > arrayLength)
{
threadGroupSize = arrayLength;
}
MTLSize threadgroupSize = MTLSizeMake(threadGroupSize, 1, 1);
应用程序向管道状态对象请求最大可能的线程组，并在该大小大于数据集大小时收缩它。该属性给出了线程组中允许的最大线程数，这取决于用于创建管道状态对象的函数的复杂性。maxTotalThreadsPerThreadgroup
编写Compute命令以执行线程
最后，编码命令以调度线程网格。
[computeEncoder dispatchThreads:gridSize
threadsPerThreadgroup:threadgroupSize];
当GPU执行此命令时，它使用您先前设置的状态和命令的参数来分派线程以执行计算。
您可以使用编码器执行相同的步骤，将多个计算命令编码到计算传递中，而不执行任何冗余步骤。例如，您可以设置一次管道状态对象，然后为要处理的每个缓冲区集合设置参数并编码命令。
结束计算通行证
如果没有其他命令要添加到计算传递，则结束编码过程以关闭计算传递。
[computeEncoder endEncoding];
提交命令缓冲区以执行其命令
[commandBuffer commit];
命令队列创建了命令缓冲区，因此提交缓冲区始终将其放在该队列上。提交命令缓冲区后，Metal异步准备执行命令，然后调度命令缓冲区以在GPU上执行。GPU执行命令缓冲区中的所有命令后，Metal将命令缓冲区标记为完成。
等待计算完成
当GPU处理您的命令时，您的应用程序可以执行其他工作。此示例不需要执行任何其他工作，因此只需等待命令缓冲区完成。
[commandBuffer waitUntilCompleted];
或者，要在Metal处理完所有命令时收到通知，请在命令缓冲区（）中添加完成处理程序，或通过读取其属性来检查命令缓冲区的状态
从缓冲区读取结果
命令缓冲区完成后，GPU的计算存储在输出缓冲区中，Metal执行任何必要的步骤以确保CPU可以看到它们。在真实的应用程序中，您将从缓冲区读取结果并对其执行某些操作，例如在屏幕上显示结果或将其写入文件。由于计算仅用于说明创建Metal应用程序的过程，因此示例将读取存储在输出缓冲区中的值并进行测试以确保CPU和GPU计算出相同的结果。
- (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]);
}
}
