咱们许多人多少都知道一些、或者听说过 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)是什么?
烘托管线,实际上指的是一堆原始图形数据途经一个输送管道,期间经过各种变化处理最终呈现在屏幕的进程。通常情况下,烘托管线有三个阶段,其间光栅化阶段不行编程,其他两个能够。
- 极点上色器(vertex stage),接收一组极点数据数组,来限定显现/处理区域。
- 光栅化阶段(rasterization stage),在光栅化阶段,确认哪些像素坐落鸿沟,裁剪超出鸿沟的像素。
- 片段上色器(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 上运转。
- 编写 MSL 函数。
- 首要找到一个可用的 GPU。
- 创立 pipeline。
- 创立 GPU 可拜访的数据目标。
- 创立指令缓冲区,向其间写入指令。
- 将缓冲区提交到指令行列。
写一个 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/