我们很多人多少都知道一些、或者听说过 Metal。Metal 是 Apple 公司开发的一套兼顾图形与计算功能，面向底层、低开销的硬件加速。其类似于将 OpenGL 与 OpenCL 的功能集成到了同一个 API 上，最初支持的系统是 iOS 8， Metal 使得 iOS 可以实现其他平台的类似功能，例如 Khronos Group 的跨平台 Vulkan 与 Microsoft Windows 上的 Direct3D 12。

Metal也通过引入计算着色器来进一步提高 GPGPU（图形处理器通用计算） 编程的能力。

Metal 使用一种基于 C++11 的新着色语言，其实现借助了 Clang 和 LLVM。

历史

• 2014 年 6 月 2 日，Metal 开始支持 iOS 设备（仅支持 Apple A7 或更新款处理器的 iPhone、iPad）。
• 2015 年 6 月 8 日，Metal 开始支持运行 OS X El Capitan 的 Mac 设备（仅 2012 年中或更新款机种）。
• 2017 年 6 月 5 日，Apple 于 WWDC 宣布了 Metal 的第二个版本，支持 macOS High Sierra、iOS 11 和 tvOS 11。 Metal 2 不是 Metal 的独立 API，并且由需要的硬体支援。 Metal 2 在 Xcode 中实现了更高效的分析和调试，加速了机器学习、降低了 CPU 工作负载、支持 macOS 上的虚拟实境以及 Apple A11 处理器的特性。
• 2019年 6 月 3 日，Metal API 更新到第三个版本，支持 macOS Catalina、iOS 13 和 iPadOS 13。
• 2020 年的苹果全球开发者大会（WWDC）上，苹果宣布将 Mac 迁移到 Apple Silicon。使用 Apple Silicon 的 Mac 将使用 Apple GPU，支持之前在 macOS 和 iOS 上实现的特色功能，并将能够利用为 Apple GPU 架构所定制的基于图块的延迟渲染（TBDR）功能。

Metal 是什么？（来自 Apple）

使用 GPU 处理器渲染高级 3D 图形和并行计算数据。

我们的应用程序通过 Metal 框架可以直接访问设备的图形处理单元 (GPU)。 借助 Metal，应用程序可以利用 GPU 快速渲染复杂场景和并行运行计算任务。 例如，以下这些类的应用可以使用 Metal 来最大化它们的性能：

• 游戏复杂 3D 环境渲染。
• 视频处理应用程序，例如 Final Cut Pro。
• 科学研究应用程序，分析和处理大型数据集的。

Metal 还可以与其他框架配合使用。 例如，MetalFX 可以比原生渲染用更少的时间完成渲染（DLSS），而 MetalKit 简化了在屏幕上显示您的 Metal 内容的任务。 Metal Performance Shaders 框架可以利用每个独特的 GPU 硬件，提供了一系列的计算优化、渲染着色器库。

很多高级苹果框架都使用了 Metal，包括 RealityKit、SceneKit、SpriteKit 和 Core Image。这些高级框架帮我们实现了 GPU 编程细节。当然，我们为了更好的性能也可以通过编写自己的自定义 Metal 和shader 代码来获得更好的性能。

MetalFX 官方演示

Metal 是什么？简言之就是 Metal 提供了一组简明的 API，使得我们开发者拥有了直接和 GPU 对话的能力。

Metal 优点

• 最佳的 GPU 性能、更低的 CPU 开销 （为啥是提升 CPU 的而不是 GPU。因为 Metal 是通过 CPU 生成 GPU Commands 再交给 GPU 执行，Metal 相比 OpenGL 大大提升了 GPU Commands 转换的速度）。
• 最大程度的保证 CPU、GPU 的协作能力，使它们两人能同时处于忙碌状态，而不是会有一个是闲置的。
• 简单易上手的 API。
• Metal 在构建时完成源码编译，在运行时作为一个 library 加载，减少编译 Metal 造成的一大笔开销。

CPU 和 GPU

CPU 和 GPU 的结构如下图，我们可以很直观看到，GPU 有很多的 ALU，GPU 可以通过在各个计算单元之间分配工作负载，共同处理大量数据，GPU是专为执行复杂的数学和几何计算而设计的

• ALU（Arithmetic and Logic Unit），即计算单元。
• Cache，进行高速数据交换的存储器。
• DRAM（Dynamic Random Access Memory），即动态随机存取存储器，最为常见的系统内存。

CPU 到 GPU 计算的示例

youtu.be/-P28LKWTzrI

渲染管线（Render Pipeline）是什么？

渲染管线，实际上指的是一堆原始图形数据途经一个输送管道，期间经过各种变化处理最终出现在屏幕的过程。通常情况下，渲染管线有三个阶段，其中光栅化阶段不可编程，其他两个可以。

1. 顶点着色器（vertex stage），接收一组顶点数据数组，来限定显示／处理区域。
2. 光栅化阶段（rasterization stage），在光栅化阶段，确定哪些像素位于边界，裁剪超出边界的像素。
3. 片段着色器（fragment stage），计算每一个像素最终的颜色值。

Core Animation、Core Graphic、Metal 的区别

在架构设计上，Metal 和 Core Graphics 属于同一级，最接近 Graphics Hardware，不同的是 Metal 使用的是GPU 的渲染能力，Core Graphics 使用的是 CPU 的渲染能力。

而 Core Animation 处于更高的一级，Core Animation 通过使用 Metal 和 Core Graphics 去完成绘制。

但是封装层级越高，往往意味着性能越差，功能越少。因此，为了追求更好的性能和实现更强大的功能，我们可以绕过 Core Animation，直接调用 Metal 和 Core Graphics 渲染、绘制图像。但是也说了，GPU 在图像处理上更有优势，因此对于图像处理需求尽量使用 Metal。

利用 GPU 进行计算

在此示例中，我们将了解所有 Metal 应用程序都要使用的基本任务。 我们会将用 C 编写的简单函数转换为 Metal 着色语言 (MSL)，使得它可以在 GPU 上运行。

1. 编写 MSL 函数。
2. 首先找到一个可用的 GPU。
3. 创建 pipeline。
4. 创建 GPU 可访问的数据对象。
5. 创建命令缓冲区，向其中写入命令。
6. 将缓冲区提交到命令队列。

写一个 GPU 函数来执行计算

为了展示 GPU 编程，我们将两个数组的对应元素添加在一起，将结果写入第三个数组作为本示例计算函数。示例 1：展示了在 CPU 上执行此计算的函数，用 C 语言编写。它遍历索引，每次循环迭代计算一个值。

示例 1：C 语言实现的数组（Array）加法

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 是 C++ 的一个变体，专为 GPU 编程而设计。在 Metal 中，运行在 GPU 上的代码被称为着色器，这个名称的由来是有历史原因的，它最初被用于计算 3D 图形中的颜色。示例 2： 展示了 MSL 中的着色器，它实现与示例 1 相同的计算。示例代码在 add.metal 文件中定义了这个函数。Xcode 在应用程序目标中构建所有 . Metal 文件，并创建一个默认的 Metal 库，并打包到我们的应用程序中。

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 版本中有一些重要的区别。

首先，函数添加 kernel 关键字，声明函数为:

• 公开 GPU 函数。公开函数只有你自己的应用能访问到。公开函数也不能被其他 shader 函数调用。
• 计算函数(也称为计算内核)，它使用线程网格执行并行计算。

更多声明，请参阅 Using a Render Pipeline to Render Primitives，了解用于声明公开 GPU 函数的其他函数关键字。

add_arrays 函数用 device 关键字声明了它的三个参数，表示这些指针位于设备地址空间中。MSL 为内存定义了几个不相交的地址空间。在 MSL 中声明指针时，必须提供一个关键字来声明其地址空间。本例使用设备地址空间来声明持有内存，以保证 GPU 可以读取和写入。

示例 2 删除了 for 循环，因为该函数现在将由计算网格中的多个线程调用。这个示例创建了一个与数组尺寸完全匹配的一维线程网格，这样数组中的每个元素都由不同的线程计算。

为了替换先前由 for 循环提供的索引，该函数接受一个新的索引参数，该参数使用另一个 MSL 关键字 thread_position_in_grid，该关键字使用 C++ 属性语法指定。这个关键字声明 Metal 应该为每个线程计算一个唯一的索引，并在这个参数中传递该索引。因为 add_arrays 使用一维网格，所以索引被定义为一个标量整数。

找到一个可用的 GPU

在你的应用中，MTLDevice 对象是 GPU 的一个抽象，你可以用它来和 GPU 通信。Metal为每个GPU创建一个 MTLDevice。通过调用 MTLCreateSystemDefaultDevice() 获得默认设备对象。在 macOS 中，Mac 可以有多个 GPU, Metal 选择其中一个 GPU 作为默认值并返回该 GPU 的设备对象。在 macOS 中，Metal 提供了其他 API，我们可以使用这些 API 来检索所有设备对象，本示例仅使用默认值。

id<MTLDevice> device = MTLCreateSystemDefaultDevice();

初始化 Metal 对象

Metal 是一个 GPU 相关的实体的集合，如编译着色器，内存缓冲区和纹理。要创建这些特定 GPU 的对象，我们可以调用 MTLDevice 上的方法，或者调用由 MTLDevice 创建的对象上的方法。由设备对象直接或间接创建的所有对象只能用于该设备对象。使用多个 GPU 的应用程序将使用多个设备对象，并为每个设备创建类似的 Metal 对象层次结构。

示例代码使用自定义 MetalAdder 类来管理与 GPU 通信所需的对象。类的初始化器创建这些对象并将它们存储在其属性中。应用程序创建该类的一个实例，传入 Metal 设备对象以用于创建次要对象。MetalAdder 对象在完成执行之前保持对 Metal 对象的强引用。

MetalAdder* adder = [[MetalAdder alloc] initWithDevice:device];

在 Metal 中，开销很大的任务建议只在开始的时候初始化一次，结果可以被保留下来并以低廉的成本使用。

获取 Metal 函数引用

初始化器做的第一件事是加载函数并准备它在 GPU 上运行。我们构建应用程序时，Xcode 编译 add_arrays 函数并将其添加到默认的 Metal 库中，并嵌入到应用程序中。我们可以使用MTLLibrary 和 MTLFunction 对象来获取有关 Metal 库和其中包含的函数的信息。要获得一个代表 add_arrays 函数的对象，通过 MTLDevice 为默认库创建一个 MTLLibrary 对象，然后请求库创建一个代表 shader 函数的 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 Pipeline

函数对象是 MSL 函数的代理，但它不是可执行代码。通过创建管道将函数转换为可执行代码。流水线指定了 GPU 为完成特定任务所执行的步骤。在 Metal 中，Pipeline State 代表 Pipeline 。因为这个示例使用了一个计算函数，所以应用程序创建了一个 MTLComputePipelineState 对象。

_mAddFunctionPSO = [_mDevice newComputePipelineStateWithFunction: addFunction error:&error];

计算 pipeline 运行单个计算函数，在运行函数之前操作输入数据，然后操作输出数据。

当你创建一个 Pipeline State 对象，设备对象为我们选定的 GPU 完成函数编译。这个示例同步创建管道状态对象，并将其直接返回给应用程序。因为编译确实需要一段时间，所以要避免在性能敏感的代码中同步创建管道状态对象。

创建命令队列

要将工作发送到 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是否可以访问它。示例应用程序使用共享内存(storagemodeshare)， CPU和GPU都可以访问。

为了用随机数据填充缓冲区，应用程序获取一个指向缓冲区内存的指针，并在 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];

要对命令进行编码，需要对编码器进行一系列方法调用。有些方法设置状态信息，如管道状态对象(PSO)或传递给管道的参数。在进行这些状态更改之后，将编码一个命令来执行管道。编码器将所有状态更改和命令参数写入命令缓冲区。

设置 Pipeline State 和参数数据

设置要执行命令的管道的管道状态对象。然后为管道需要发送到 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 表示该命令将从缓冲区的开头访问数据。另外，我们可以使用一个缓冲区来存储多个参数，并为每个参数指定偏移量。

我们不需要为 index 参数指定任何数据，因为 add_arrays 函数将其值定义为由 GPU 提供。

指定线程组数量

接下来，决定要创建多少线程以及如何组织这些线程。Metal 可以创建一维、二维或三维网格。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);

应用程序向管道状态对象询问最大的线程组，如果该线程组的大小大于数据集的大小，则收缩线程组。maxTotalThreadsPerThreadgroup属性给出线程组中允许的最大线程数，这取决于用于创建管道状态对象的函数的复杂程度。

编码计算命令提交到执行线程

最后，对命令进行编码以分派线程网格。

[computeEncoder dispatchThreads:gridSize
          threadsPerThreadgroup:threadgroupSize];

GPU 执行此命令时，将使用我们之前设置的状态和命令的参数来调度线程来执行计算。

我们可以按照相同的步骤使用编码器将多个计算命令编码到计算通道中，而无需执行任何冗余步骤。例如，您可以设置一次管道状态对象，然后为要处理的每个缓冲区集合设置参数并编码命令。

结束计算编码

当没有更多命令要添加到计算通道时，可以结束编码过程以关闭计算通道。

[computeEncoder endEncoding];

提交到命名缓冲区并执行

通过将命令缓冲区提交到队列来运行命令缓冲区中的命令。

[commandBuffer commit];

命令队列创建了命令缓冲区，因此提交缓冲区总是将其放在该队列中。在提交命令缓冲区后，Metal 会异步准备要执行的命令，然后调度命令缓冲区在 GPU 上执行。当 GPU 执行完命令缓冲区中的所有命令后，Metal 将命令缓冲区标记为完成。

等待计算完成

当 GPU 处理我们的命令时，我们可以做其他工作。示例代码不需要做任何额外的工作，我们只需要等待命令缓冲区执行完成。

[commandBuffer waitUntilCompleted];

除此之外，可以监听 Metal 处理完所有命令时到通知，我们可以向命令缓冲区添加一个完成处理程序(addCompletedHandler(_:)))，或者通过读取命令缓冲区的status 属性来检查命令缓冲区的状态。

从缓冲区读取结果

在命令缓冲区完成后，GPU 的计算被存储在输出缓冲区中，Metal 执行任何必要的步骤以确保 CPU 可以看到它们。最后我们将从缓冲区读取结果并对其进行处理，校验 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");
}

参考

---

• developer.apple.com/documentati…
• xiaoye220.github.io/Metal-1/