Metal: 使用图形处理器渲染高级三维图形并执行数据并行计算
编写一个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上，需要通过MSL（Metal Shading Language），MSL是为GPU编程而设计的C++变体，在Metal API 中在GPU上运行的代码被称为着色器，因为在历史上他们是第一次被用来运算3D绘图上颜色。在表2中显示MSL上的一个执行和表1一样的运算的着色器。
列表2，数组的相加，SML编写
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版本中有一些重要的不同点。我们来看一下列表2.
首先，函数添加kernel关键字，该关键字声明函数为：
1.一个公共GPU函数，公共函数是在APP上能看见的唯一的函数， 公共函数也不能被其他的着色器函数调用。
2.一个运算函数（也被称之为运算内核），是通过线程网络进行平行运算的函数。
阅读使用渲染管道渲染基本体去学习其他更多的公共绘图函数的函数关键字。
add_arrays函数通过设备关键字声明了它的三个属性，这说明这些属性的指针在设备的地址空间里。MSL为内存定义了几个不连续的地址空间。无论何时， 你在MSL上定义一个指针，你必须提供一个关键字来定义它的地址空间。使用设备地址空间声明GPU可以读写的持久内存。
列表2 移除了列表1中的for循环，因为这个函数现在在运算网络中的多个线程调用，此示例创建与数组维度完全匹配的线程的一维网格，因此数组中的每个条目都由不同的线程计算。为了替换for循环提供的上一个索引，通过使用C++属性语法置顶的另一个MSL关键字（thread_position_in_grid）给函数设置一个新的索引属性，此关键字声明Metal 应该为每一个线程计算一个唯一的索引，并在该参数中传递该索引。因为add_arrays用的是一维网格，这个索引会被定义为标量整数，即便循环被移除，清单1和清单2使用同一行代码将这两个数字相加。如果你想从C和C++中转换一样的代码到SML，用同样的方法把循环逻辑替换为网格。
在应用程序中，MTLDevice对象是GPU的一个抽象,用它可以和GPU交互，Metal 为每一个GPU创建一个MTLDevice，你可以调用 MTLCreateSystemDefaultDevice()获得默认device 对象。在macOS 中，Mac 可以有多个GPU，Metal 选择其中一个作为默认GPU并返回他的device 对象。在macOS中，Metal还提供了其他的APIs用来检索所有的device 对象，但下面的例子只用于默认的对象。
id<MTLDevice> device = MTLCreateSystemDefaultDevice();
初始化Metal对象
初始化做的第一件事儿是加载函数并准备它在GPU 上运行，当你创建应用，Xcode 编译add_arrays函数并且把它添加到嵌入到应用里面的metal库中。你可以用MTLLibrary 和 MTLFunction 对象来获取metal库和函数中的相关信息。为了获取一个表示add_arrays的函数，通过 MTLDevice为默认库创建一个MTLLibrary对象并向库请求一个MTLFunction对象来表示着色器函数。
- (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;
}
准备一个metal管道
函数对象是MSL函数的一个代理，但是它不是一个可执行的代码， 通过创建一个管道把函数转换成可执行代码。管道指定了GPU完成特定任务所执行的步骤。在Metal 中，管道由管道状态对象表示。因为这个例子用了一个计算函数，应用创建一个MTLComputePipelineState对象。
_mAddFunctionPSO = [_mDevice newComputePipelineStateWithFunction: addFunction error:&error];
一个计算管道运行一个简单的计算函数，在运行函数之前可选地操作输入数据，以及在运行函数之后操作输出数据。
当创建一个管道状态对象，device对象完成这个特定GPU的函数编译。此示例同步创建管道状态对象，并将其直接返回给应用程序。因为编译确实需要一段时间，所以避免在对性能敏感的代码中同步创建管道状态对象。
提示
到目前为止您看到的代码中Metal返回的所有对象都是作为符合协议的对象返回的。Metal使用协议来抽象底层实现类来定义大多数特定于gpu的对象，这些实现类对于不同的gpu可能会有所不同。Metal使用类定义了与gpu无关的对象，任何给定的Metal协议的参考文档都明确说明了您是否可以在应用程序中实现该协议。
创建命令队列
为了把工作发送到GPU，你需要创建一个命令队列。metal使用命令队列来调度命令，通过向MTLDevice请求一个命令队列来创建一个命令队列。
_mCommandQueue = [_mDevice newCommandQueue];
创建数据缓冲区并加载数据
初始化基本metal 对象之后，为GPU 加载执行的数据，这个任务对性能的影响较小，但在应用启动的早期这样做仍然很有用。
GPU有它自己的专用内存，或者它可以与操作系统共享内存。Metal和操作系统内核需要执行额外的工作，以便让您在内存中存储数据，并使这些数据可供GPU使用。Metal使用资源对象抽象了这个内存管理。MTLResource是GPU在运行命令时可以访问的内存分配的资源。通过 MTLDevice来为GPU创建资源。
示例应用创建三个缓冲区并用随机数据填充前两个，第三个缓冲区是add_arrays存储结果的地方。
_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是否可以访问它。示例应用程序使用共享内存storageModeShared，CPU和GPU都可以访问该内存。
为了用随机数据填充缓冲区，应用程序从缓冲区内存获取一个指针并在CPU上给它写入数据。在列表2中的add_array函数声明了它的浮点数类型的数组参数。所以你需要提供一样格式的缓冲区。
- (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];
为了编码一个命令，对编码器进行一系列方法调用。有些方法设置状态信息，如管道状态对象或传递给管道的参数。改变这些状态后， 你可以对命令进行编码以执行管道。编码器将所有状态更改和命令参数写入命令缓冲区。
设置管道状态和参数数据
设置要执行命令的管道的管道状态对象，然后为管道需要发送到add_arrays函数的任何参数设置数据。对于这个管道，这意味着提供对三个缓冲区的引用。Metal自动按照清单2中函数声明中参数出现的顺序为缓冲区参数分配索引，从0开始。使用相同的索引提供参数。
[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表示命令将从缓冲区开始访问数据。不管怎么样，你可以使用一个缓冲区存储多个参数，为每个参数指定偏移量。
你没有为索引参数提供任何参数，因为add_arrays函数定义它的值由GPU提供。
指定线程数和组织
接下来，决定要创建多少线程以及如何组织这些线程。Metal 可以创建1维，2维或3维网格。add_arrays函数使用一维数组，因此该示例创建了一个大小为（dataSize x 1 x 1）的一维网格，Metal从中生成0到dataSize-1之间的索引。
MTLSize gridSize = MTLSizeMake(arrayLength, 1, 1);
指定线程组大小
Metal 将网格细分为更小的网格，成为线程组。每个线程组单独运算。Metal 可以将线程组分配给GPU上的不同的处理元素以加快处理速度。您还需要确定为命令创建线程组的大小。
NSUInteger threadGroupSize = _mAddFunctionPSO.maxTotalThreadsPerThreadgroup;
if (threadGroupSize > arrayLength)
{
threadGroupSize = arrayLength;
}
MTLSize threadgroupSize = MTLSizeMake(threadGroupSize, 1, 1);
应用程序向管道状态对象请求最大可能的线程组，如果该大小大于数据集的大小，则会收缩该线程组。 maxTotalThreadsPerThreadgrou属性给一个线程组允许的最大的值，它根据用于创建管道状态对象的函数的复杂性的不同而不同。
编码运算命令以执行线程
最后，对命令进行编码以调度线程网格。
[computeEncoder dispatchThreads:gridSize
threadsPerThreadgroup:threadgroupSize];
当GPU执行此命令时，它使用你先前设置的状态和命令的参数来调度线程来执行计算。你可以按照相同的步骤使用编码器将多个计算命令编码到计算过程中，而无需执行任何冗余步骤。例如，可以设置管道状态对象一次，然后为要处理的每个缓冲区集合设置参数并编码一个命令。
结束计算过程
当没有更多的命令添加到计算过程时，结束编码过程以结束计算过程。
[computeEncoder endEncoding];
提交命令缓冲区以执行其命令
通过将命令缓冲区提交到队列来执行命令缓冲区中的命令。
[commandBuffer commit];
命令队列创建了命令缓冲区，因此提交缓冲区时总是将其放在该队列上。提交命令缓冲区后，Metal 异步的准备要执行的命令，然后调度命令缓冲区在GPU上执行。在GPU执行完命令缓冲区的所有命令后，Metal将命令缓冲区标记为已完成。
等待运算完成
当GPU处理命令时，应用可以做其他的工作。这个样本不需要做任何额外的工作，所以它只需要等待 命令缓冲区完成。
[commandBuffer waitUntilCompleted];
或者，为了在Metal 处理完所有命令时得到通知，向命令缓冲区添加一个完成处理程序(addCompletedHandler(_:))，或者通过读取命令缓冲区的状态status属性来检查它的状态。
从缓冲区读取结果
命令缓冲区完成后，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++)
{
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");
}
网友评论