本文翻译自苹果Metal的官方文档:Performing Calculations on a GPU,如果哪里翻译有误,请大家指出,我会及时更正。欢迎转载,但禁止用于商业用途。谢谢!
概述
在这个例程中,你将学习到在所有Metal程序中所用到的基本任务。你将看到如何把一个简单的c语言程序转换成Metal着色器语言(MSL),并使其能够在GPU上运行。你会找到一个GPU,通过创建管线准备使MSL函数在GPU上运行,并构造GPU能够访问的数据对象。针对你的数据对象去运行管线,首先创建一个命令缓存区,将命令写入其中,然后将缓存区提交到命令队列,最后Metal将命令发送到GPU,并执行它们。
编写一个GPU函数来执行计算
为讲解GPU编程,给出一个两数组对应位求和,并把结果放进第三个数组的示例程序。示例1为在cpu上运行上述运算的C程序,通过遍历索引index,在每次循环中计算result。
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的着色器语言(MSL)。MSL为GPU编程而设计的c++变体。其中,运行在GPU的代码被称为着色器(shader),这是由于这类GPU程序曾首先被应用于计算3D图像中的颜色。示例2展示了与示例1进行同样运算的MSL shader。在官方给出的示例代码中,这个函数定义在add.metal。使用Xcode构建并创建一个默认的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];
}
示例2与示例1相似,当在MSL版本中,有一些重要的区别。仔细观察示例2.
首先,这个函数加了一个kernel关键字,使这个函数被声明为:
- 公共(public)GPU函数,程序在调用这个Metal库时,public是仅可见的函数,同时public函数不能被其他shader程序调用。
- 计算函数(也称为计算kernel),它通过线程网格(gird)完成并行计算。
如需了解其他用于声明public图形函数的关键字,请参阅使用渲染管道渲染图元。
add_arrays 函数用 device 关键字声明了它的三个参数,这表示这些数组指针位于Device地址空间中。MSL定义了若干不相邻的内存区域,每当在MSL中声明一个指针,必须使用一个关键字来声明它的内存地址。使用device修饰的地址空间,代表GPU可以对其进行持续的读写。
示例2移除了示例1的for循环,因为是被gird中的多线程调用的,这个示例创建了一个与数组维度一致的一维线程网格,因此数组中每一个数值求和是不同线程完成计算的。
同时示例2使用了一个新的index索引变量替代之前在for循环中的索引,index通过c++的attribute被赋值为另一个MSL关键词:thread_position_in_grid,这个关键字表明 Metal 会为每个线程计算一个唯一的索引,并将索引传递给变量。因为add_arrays使用了一维gird,因此index定义为整数(uint)。尽管for循环被移除,示例1和示例2使用相同的方式完成求和。如果想将类似的代码从c/c++转换为MSL,可以使用相同的方式循环逻辑的替换。
检索GPU
在应用程序中,MTLDevice代表一个GPU的简单抽象(Metal设备对象),可通过它完成与GPU的通信。Metal为每一个GPU创建一个MTLDevice,并可通过调用MTLCreateSystemDefaultDevice()获取默认的设备对象。在 macOS 中,Mac 可以有多个 GPU,Metal 会选择其中一个 GPU 作为默认值并返回该 GPU 的MTLDevice,Metal也提供了用于检索所有GPU的其他API,但此示例仅使用默认值。
id<MTLDevice> device = MTLCreateSystemDefaultDevice();
初始化Metal对象
Metal将其他与GPU相关的工作(如编译shader、内存缓存区、图像纹理)表示为不同对象。要创建这些特定的GPU对象。可以通过调用MTLDevice和派生于MTLDevice对象的方法创建这些特定的GPU对象。由GPU的MTLDevice对象直接或间接创建的所有对象仅可用于对应的GPU,对于使用多个GPU的应用程序,Metal会创建多个GPU对象,并为每个对象创建类似的Metal层次结构。
示例程序中使用一个自定义的类MetalAdder来管理需要与GPU进行通信的对象,在类的初始化中会创建这些对象,并存储在其属性中。在程序中会创建一个该类的示例,传入用于创建辅助对象的 Metal 设备对象。这个MetalAdder对象会在生命周期内保持对Metal对象的强引用。
MetalAdder* adder = [[MetalAdder alloc] initWithDevice:device];
在Metal中,高消耗的初始化任务可以一次运行之后,将结果保留并在低消耗情况下被使用。很少情况下需要在性能要求严格的代码中执行此类任务。
调用Metal函数
初始化首先加载这些函数,同时准备在GPU上运行他们。当在build示例程序时,XCode会编译add_arrays函数到一个默认的Metal库,后续应用层会完成对其的调用。使用MTLLibray和MTLFunction的对象去获取Metal库和库中所包含的函数。为获取代表add_array的对象首先需要让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渲染管线
上面的函数对象addFunction代表已完成的MSL函数,但它并不是能够直接运行的,需要通过创建一个管线(pipline),使其能够被执行。pipline会让GPU明确完成一个特定任务所需要执行的步骤。在Metal,pipline会被一个管线状态对象所表示,因为在例程中,两数求和是一个计算函数,因此使用MTLComputePiplineState完成管线对象的创建。
_mAddFunctionPSO = [_mDevice newComputePipelineStateWithFunction: addFunction error:&error];
计算pipline会运行单个计算函数,可选地在运行函数之前处理输入数据,运行之后处理输出数据。
当创建管线状态对象时,设备对象会为特定的GPU编译好需要执行的函数。在示例中,会同步创建管道状态对象,并将其直接返回给应用程序。 因为编译确实需要一段时间,所以避免在性能敏感代码中同步创建管道状态对象。
原文注:到目前为止,在代码中看到的所有 Metal 返回的对象都是作为符合协议的对象返回的。 Metal 使用协议定义了大多数GPU的特有对象,以完成底层实现类的抽象,这些实现类可能因不同的 GPU 而异。 Metal 使用这些类定义与GPU所独有的对象。 对于是否可以在应用程序中实现该协议可以参阅Metal的参考文档。
创建一个指令队列(Command Queue)
向GPU发送工作任务,需要使用指令队列(command queue)实现,同时Metal会使用指令队列,完成对指令的调度。通过MTLDevice完成指令队列的创建。
_mCommandQueue = [_mDevice newCommandQueue];
创建数据缓存区并加载数据
在创建好基本的Metal对象之后,需要加在数据以供GPU执行,此任务对性能的要求不高,但最好也在应用程序启动时执行,避免性能损耗。
GPU可以拥有自己的专用内存,也可以与操作系统共享内存。Metal和操作系统内核需要执行额外的工作才能将数据存储在内存中,并使这些数据可供GPU使用。Metal使用资源类(MTLResource)完成内存资源的管理。resource是GPU在运行命令时可以访问的内存区域。可使用MTLDevice完成对其GPU的resource的创建。
在示例程序中,创建了三个MTLBuffer,其中前两个使用随机数填充,第三个用来存放add_array的结果。
_mBufferA = [_mDevice newBufferWithLength:bufferSize options:MTLResourceStorageModeShared];
_mBufferB = [_mDevice newBufferWithLength:bufferSize options:MTLResourceStorageModeShared];
_mBufferResult = [_mDevice newBufferWithLength:bufferSize options:MTLResourceStorageModeShared];
[self generateRandomFloatData:_mBufferA];
[self generateRandomFloatData:_mBufferB];
在本示例中,resource为这些MTLBuffer对象,它们是没有预定义格式的内存分配。Metal 将每个缓冲区作为不透明的字节集合进行管理。 但是在shader中使用缓冲区时会指定数据格式。 这意味着shader和应用程序需要就来回传递的任何数据的格式进行统一。
当分配一个buffer时,需要提供一种存储形式来确定它的一些性能特征,同时需明确CPU/GPU是否可以访问它。在示例程序中,使用共享内存(storgeModeShared)来保证CPU和GPU能够访问。
为向buffer内填随机数,应用程序获取了buffer的指针,并使用cpu向其中写数据。示例2中的arr_arrays函数声明了它的参数为float型的数组指针,所以需要保证buffer内的数据与其类型一致。
- (void) generateRandomFloatData: (id<MTLBuffer>) buffer
{
float* dataPtr = buffer.contents;
for (unsigned long index = 0; index < arrayLength; index++)
{
dataPtr[index] = (float)rand()/(float)(RAND_MAX);
}
}
创建一个指令缓存区(Command Buffer)
使用指令队列(_mCommandQueue)对象创建一个指令缓存区。
id<MTLCommandBuffer> commandBuffer = [_mCommandQueue commandBuffer];
创建一个指令编码器(Command Encoder)
为了向指令缓存中写指令,需要使用Command Encoder来处理编码时的特定指令。此示例创建一个计算命令编码器,该编码器对计算通道进行编码。 计算通道包含执行计算管道的命令列表。 每个计算命令都会使 GPU 创建一个线程网格以在 GPU 上执行。
id<MTLComputeCommandEncoder> computeEncoder = [commandBuffer computeCommandEncoder];
要对指令进行编码,需要对编码器进行一系列方法的调用,其中主要为设置状态信息,如管道状态对象 (pipeline state object,PSO) 或要传递给管道的参数。 进行这些状态更改后,命令进行编码以执行pipline。编码器会将所有状态变化和命令参数写入命令缓冲区。
设置管线状态和参数信息
首先需要设置对应pipline的状态,即PSO,然后针对需要发送到add_arrays函数的每一个参数进行配置。对于此pipline,需要对三个buffer进行配置。 Metal 按照参数出现在示例2中的顺序自动为buffer参数分配索引,从 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];
同时也需要制定参数的偏移量(offset),当偏移量为0时,表明command会从buffer的起始位置获取参数。但也可以使用一个buffer存储多个参数,为每一个参数指定一个独立的offset。
可以看出,程序中没有为add_arrays的index参数指定任何数据,这是由于该值由GPU提供。
指定线程数并组织线程
下一步,确定需要创建多少线程,并如何组织这些线程。Metal能够创建一维、二维、三维的线程网格。由于add_arrays使用的是一维数组,所以例程中创建了一个长度为数组长度的一维网格
MTLSize gridSize = MTLSizeMake(arrayLength, 1, 1);
指定线程组(Threadgroup)大小
Metal为网格细分为更小的结构,叫做线程组(Threadgroup)。每个线程组都是单独计算的。 Metal 可以将线程组分派给 GPU 上的不同处理元素以加快处理速度。 同样也需要决定为指令创建对应尺寸的线程组。
NSUInteger threadGroupSize = _mAddFunctionPSO.maxTotalThreadsPerThreadgroup;
if (threadGroupSize > arrayLength)
{
threadGroupSize = arrayLength;
}
MTLSize threadgroupSize = MTLSizeMake(threadGroupSize, 1, 1);
在例程中,使用PSO获取了最大的线程组尺寸,如果这个长度大于参数的本身长度,则需要调整threadGroupSize为对应的大小。
maxTotalThreadsPerThreadgroup 属性给出线程组中允许的最大线程数,这取决于用于创建管道状态对象的函数复杂性。
编码计算指令并在线程中执行
最后,对指令进行编码去分配线程网格。
[computeEncoder dispatchThreads:gridSize
threadsPerThreadgroup:threadgroupSize];
当GPU运行这个指令,会使用先前设置好的状态和命令的参数来调度线程来执行计算。 可以使用编码器按照相同的步骤将多个计算命令编码到计算通道中,而无需执行任何冗余步骤。 例如,可以设置管道状态对象一次后,设置参数为要处理的每个缓冲区集合编码一个命令。
结束计算通道
当没有更多的指令进入计算通道时,需要结束编码关闭计算通道。
[computeEncoder endEncoding];
提交命令缓冲区以执行其命令
指令缓冲(command buffer)提交命令到指令队列(command queue),然后queue来运行命令缓冲区中的命令。
[commandBuffer commit];
指令队列创建了指令缓存,因此提交缓存总是将其放入队列中。在提交指令缓存后,Metal会进行预处理操作,然后安排command buffer在GPU上执行。GPU在执行完缓存中的command之后,Metal会将缓存区标记成已完成。
等待计算完成
在GPU处理command的时候,应用程序可以完成其他的工作,在示例程序中,应用层不需要处理额外的工作,所以只需要简单地等待command buffer完成任务。
[commandBuffer waitUntilCompleted];
同样的,也可在Metal完成command的时候,向command buffer添加一个完成的句柄(handler),也可以读取command buffer的状态属性,获取command是否完成。
从缓存中读取结果
在command buffer完成之后,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");
}