一、在GPU上执行运算

本文Demo

环境：

mac os 10.14.5

xcode 10.3

此系列文章源自官方案例，详情至此处

专用名词虽有汉字翻译，但会保留原有英文形式名词。

概述

在此示例中，会学习在所有 Metal apps 中使用到的基本要素：

a）把用 C 写的简单函数转化成 Metal Shading Language（MSL），因此可以在 GPU 上运行

b）找到 GPU

c）通过创建管道准备在 GPU 上运行 MSL 函数

d）创建 GPU 可访问的内存分配以保存数据

e）创建命令缓冲区和编码GPU命令去控制数据

f）提交缓冲区到命令队列，使GPU执行编码命令

1.写一个GPU函数执行计算

将两个具有相同元素个数的数组相加，结果放入第三个数组。下面计算函数用 C 编写，运行在 CPU 上。

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 重写该方法。 MSL 是专为 GPU编程设计的 C++ 的变种。在 Metal中，运行在 GPU 上的代码叫做shader，因历史缘故，它们第一次在3D绘制中被用来计算颜色。

示例在 add.metal 文件中定义了用 MSL编写的 shader。Xcode会编译所有的.metal 文件，生成默认的 Metal库 导入app。

kernel void add_arrays(device const float* inA,device const float* inB,device float* result,uint index [[thread_position_in_grid]])
{// for循环换成了线程集合，在此被调用result[index] = inA[index] + inB[index];
}

由上可知：

需要添加 kernel关键字，声明此函数是：

a）公共的GPU函数。公共函数是app能够调用的唯一函数。公共函数不能被其他 shader函数调用。

b）计算函数（也称计算内核）。使用线程网格执行并行计算。

add_arrays 函数的前三个参数有 device关键字，表明指针在 device地址空间。MSL 为内存定义几个不相交的地址空间。在 MSL中，无论何时申请指针，都需要提供关键字声明地址空间。device地址空间 声明了GPU可以读取和写入的持久内存。

之所以函数移除for 循环，因为现在在网格计算中通过多线程被调用。示例创建1D的网格线程与数组的大小匹配，因此数组中的每一个元素都通过不同的线程计算。

add_arrays 函数的第四个参数，使用MSL的 thread_position_in_grid关键字，指定使用C++属性语法。这个关键字声明 Metal 应该为每个线程计算一个独特的索引，在参数中传递。由于add_arrays 函数使用1D网格，索引当作标量整数。

2.发现GPU

MTLDevice对象是GPU的精简抽象，使用它与GPU建立连接。Metal 为每个GPU创建 MTLDevice对象。通过调用 MTLCreateSystemDefaultDevice 获取默认的 device对象。在MacOS中，有多个GPUs的Mac，Metal 选择其中一个作为默认 device。

id<MTLDevice> device = MTLCreateSystemDefaultDevice();

3.初始化Metal对象

示例自定义 MetalAddr类管理与GPU连接的对象。类初始化时创建这些对象并存储在属性中。

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

在 Metal 中，费时的任务可以运行一次并且保存结果，使用时方便。在性能要求高的代码中很少使用此类任务。

4.获取Metal功能的引用

当编译app时，Xcode会编译 add_arrays 函数到默认的 Metal库中。用 MTLLibrary和 MTLFunction对象来获取 Metal库 和所包含函数的相关信息。想要获得add_arrays 函数的引用，需要 MTLDevice创建 MTLLibrary对象，然后通过 MTLLibrary获取函数引用。

- (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 projectid<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;}

5.准备Metal Pipeline（管道）

该函数对象是 MSL函数 的代理，不是可执行的代码。通过创建 pipeline将函数转变成可执行的代码。Pipeline（管道）指定GPU执行完成指定任务的步骤。在 Metal 中，用 pipeline state object（管道状态对象）代表一个Pipeline（管道）。

示例使用计算功能，因此创建 MTLComputePipelineState 对象。

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

计算管道运行单个计算函数，可选地在运行函数之前操作输入数据，然后输出数据。

当创建 pipeline state object 时，device对象将完成为指定GPU编译的功能。示例同步创建 pipeline state object，因为编译需要一些时间，因此避免在性能要求的代码中使用。

6.创建Command Queue（命令队列）

向GPU发送任务，需要命令队列。Metal 使用命令队列去安排命令。需要 MTLDevice 创建命令队列：

_mCommandQueue = [_mDevice newCommandQueue];

7.创建数据缓冲区以及加载数据

初始化基本 Metal对象 之后，加载要GPU执行的数据。此任务的性能不太重要，但是在应用程序启动的早期仍然有用。

GPU可以拥有专有内存，或者与操作系统共享内存。Metal 和操作系统内核需要执行额外的工作，将数据保存在内存中，供GPU使用。Metal 用 resource对象（MTLResource）抽象内存管理。resource（资源）是运行命令时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 将每个缓冲区作为不透明的字节集合。然而，当你在 shader 中使用缓冲区指定格式。意味着 shader 和 app 来回传递的任何数据格式需要一致。

分配缓冲区时，需要提供存储模式去确定它的性能参数以及CPU或GPU是否可以访问。示例使用了共享内存（MTLResourceStorageModeShared），即CPU和GPU都可访问。

要用随机数填充缓冲区，app获得缓冲区的内存指针，并且在CPU上写数据。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);}
}

8.创建 Command Buffer（命令缓冲区）

请求命令队列创建命令缓冲区

id<MTLCommandBuffer> commandBuffer = [_mCommandQueue commandBuffer];

9.创建 Command Encoder（命令编码器）

要将命令写入 命令缓冲区，可以使用 命令编码器来处理要编码的特定命令类型。

示例创建一个计算命令编码器，用于编码计算传递过程。计算传递包含执行计算管道的命令列表。每个计算命令都会让GPU创建一个在GPU上执行的线程网格。

id<MTLComputeCommandEncoder> computeEncoder = [commandBuffer computeCommandEncoder];

要编码一个命令，需要在编译器上调用一系列方法。一些方法设置状态信息，像 pipeline state object（PSO 管道状态对象）或传递给管道的参数。做完状态改变之后，编码命令去执行管道。编译器将所有的状态改变和命令参数写入命令缓冲区。

10.设置管道状态和参数数据

设置要执行命令的管道的管道状态对象。然后为管道需要发送到 add_arrays函数的所有参数设置数据。管道提供对三个缓冲区的引用。Metal 自动为函数的缓冲区参数提供索引，以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提供。

11.指定线程个数和架构

接下来，决定线程个数以及架构。Metal 可以创建1D、2D和3D的网格。因 add_arrays函数用1D的数组，所以示例创建1D的网格（datasize x 1 x 1），索引从 0 到 datasize -1。

MTLSize gridSize = MTLSizeMake(arrayLength, 1, 1);

12.指定线程组大小

Metal 将网格细分成更小的网格称作 线程组。每个线程组单独计算。Metal 派发线程组在GPU上去处理不同的元素，以加快处理速度。决定为命令创建多大的线程组：

NSUInteger threadGroupSize = _mAddFunctionPSO.maxTotalThreadsPerThreadgroup;
if (threadGroupSize > arrayLength)
{threadGroupSize = arrayLength;
}
MTLSize threadgroupSize = MTLSizeMake(threadGroupSize, 1, 1);

示例获取 pipeline state object（管道状态对象）的可能最大的线程组，如果大于设置数据的大小则缩小它。 MaxTotalThreadsPerThreadroup 属性获得允许最大线程组，它的变化取决于创建管道状态对象的函数的复杂性。

13.编写计算命令以执行线程

最后，编写命令派发网格线程。

[computeEncoder dispatchThreads:gridSizethreadsPerThreadgroup:threadgroupSize];

当GPU执行这条命令时，使用先前设置的状态和命令的参数派发线程以执行计算。

可以使用编码器执行相同的步骤，编写多个计算命令放入计算传递过程，而不执行冗余的步骤。例如，可能设置管道状态对象一次，然后为要处理的每个缓冲区集合设置参数以及编码命令。

14.结束计算传递过程

当没有命令添加进计算传递过程，结束编码处理进而关闭计算传递过程。

[computeEncoder endEncoding];

15.提交命令缓冲区以执行命令

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

[commandBuffer commit];

命令队列创建缓冲区，因此提交总是在队列上。提交之后，Metal 准备异步执行命令并且计划缓冲区在GPU上执行。GPU执行完缓冲区的所有命令之后，Metal 标记缓冲区结束。

16.等待计算完成

当GPU处理命令时可以做其他任务。示例没有其他任务，所以等待直到任务完成。

[commandBuffer waitUntilCompleted];

如果要在Metal处理完所有的命令后收到通知，需要添加完成处理程序（addCompletedHandler:）；或者通过 status属性获取缓冲区状态。

17.从缓冲区读取结果

命令缓冲区完成之后，GPU的计算结果存储在输入的缓冲区，Metal执行一些步骤之后确保cpu可以访问。在app中，可能想从缓冲区读取结果并且做一些处理（如在屏幕展示或者写入文件）。由于计算结果仅用作于创建 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]);}
}

个人博客：https://blog.csdn.net/Crazy_SunShine

Github：https://github.com/cxymq

个人公众号：Flutter小同学

个人网站：http://chenhui.today/

一、在GPU上执行运算相关推荐

GPU上创建目标检测Pipeline管道
GPU上创建目标检测Pipeline管道 Creating an Object Detection Pipeline for GPUs 今年3月早些时候,展示了retinanet示例,这是一个开源示例 ...
GPU上的随机森林：比Apache Spark快2000倍
作者|Aaron Richter 编译|VK 来源|Towards Data Science 随机森林是一种机器学习算法,以其鲁棒性.准确性和可扩展性而受到许多数据科学家的信赖. 该算法通过boots ...
取模、乘法和除法运算在CPU和GPU上的效率
问题: 将整数n分解为i和j,满足下面关系: n = j * idim + i 其中idim为常量. 以下为三种算法实现: 1) i = n % idim,j = (n - i) / id ...
取模乘法和除法运算在CPU和GPU上的效率
分享一下我老师大神的人工智能教程!零基础,通俗易懂!http://blog.csdn.net/jiangjunshow 也欢迎大家转载本篇文章.分享知识,造福人民,实现我们中华民族伟大复兴! 问题: ...
GPU上稀疏矩阵的基本线性代数
GPU上稀疏矩阵的基本线性代数 cuSPARSE库为稀疏矩阵提供了GPU加速的基本线性代数子例程,这些子例程的执行速度明显快于仅CPU替代方法.提供了可用于构建GPU加速求解器的功能.cuSPARSE ...
OpenCV之gpu 模块. 使用GPU加速的计算机视觉：GPU上的相似度检测(PNSR 和 SSIM)
GPU上的相似度检测(PNSR 和 SSIM) 学习目标在 OpenCV的视频输入和相似度测量教程中我们已经学习了检测两幅图像相似度的两种方法:PSNR和SSIM.正如我们所看到的,执行这些算法需 ...
遇见C++ AMP：在GPU上做并行计算
遇见C++ AMP:在GPU上做并行计算 Written by Allen Lee I see all the young believers, your target audience. I see ...
在GPU上运行，性能是NumPy的11倍，这个Python库你值得拥有
导读:NumPy是数据计算的基础,更是深度学习框架的基石.但如果直接使用NumPy计算大数据,其性能已成为一个瓶颈. 随着数据爆炸式增长,尤其是图像数据.音频数据等数据的快速增长,迫切需要突破NumP ...
在GPU上运行MATLAB程序
matlab在运行一些大型程序时会比较慢,如果你的电脑正好有一张不错的显卡,那么为什么不用显卡来加速matlab运行呢? 本文将讲解如何使用gpu来加速matlab运行程序,并总结适合gpu加速的ma ...

一、在GPU上执行运算

一、在GPU上执行运算相关推荐

最新文章

热门文章