这篇笔记主要写一下关于Metal中的Compute Shader有关的部分。
使用显卡进行计算
并行计算
显卡的优势在于并行计算(parallelism)。对于线程安全的循环,我们就可以将其从 CPU 转移到 GPU 进行计算。
例如,在 Metal 的官方文档中使用的例子:
void add_arrays(const float *inA, const float *inB,
float* result, int length)
{
for (int i = 0; i < length; i++)
{
result[i] = inA[i] + inB[i];
}
}
这个例子中,数组中的每个元素都是独立被读取的,因此不会出现的顺序影响计算结果的问题。这种循环就很适合改成并行计算,将循环中的每一个迭代分配给单独的运算单元进行计算。
使用 Metal 将上面的代码改成 GPU 代码,则会变成
kernel void add_arrays(device const float *inA, device const float *inB, device float *result, uint i [[thread_position_in_grid]])
{
result[i] = inA[i] + inB[i];
}
在上面,加粗的部分就是 Metal 所做的变化。需要注意的点有以下几点:
循环结构没了。因为现在这个add_arrays只负责一次迭代,而具体的迭代任务则是分配到了不同的计算单元中。
kernel关键字。kernel 标记一个公有的GPU函数。只有public GPU函数才可以被我们的应用看见,但即使是public的function,其他的shader function也不能调用。kernel 同时标记该函数为计算函数,也就是可以被独立线程平行计算的函数。
device关键字。device 标记出这些指针设备记录在设备地址空间的。
[[thread_poision_in_grid]]关键字。这是一种 C++ 标记属性的语法。这个关键字声明 Metal 需要为每一个线程单独计算一个不同的索引值 i,并且把这个索引值传递给这个函数作为一个实参。
寻找 GPU 设备
在 Metal#1 的初始化阶段我们已经介绍过这个概念。我们需要一个抽象来指代我们的显卡设备,否则 device 将没有意义。GPU 的抽象在 Metal 中是一个 MTLDevice。复习一下,获得一个默认 GPU 的方法是
id<MTLDevice> device = MTLCreateSystemDefaultDevice();
虽然 macOS 支持 mac 使用多个 GPU (可能是真的),Metal 会选择其中之一的 GPU 作为默认设备被这个函数返回。
Metal 也会给其它的一些相关概念提供抽象,例如 shader, buffer和纹理,将它们以对象的形式返回。我们只需要调用相关的 MTLDevice 方法即可。
获取 Metal 函数的引用
Initializer 首先会将加载这个函数,然后做好准备让它能够在 GPU 上运行。当我们编译我们写的 App 的时候, Xcode 会将这个函数(比如上面那个 add_arrays ) 添加到 Metal 的默认函数库,将其嵌入 App。然后,我们就可以使用 MTLLibrary 和 MTLFunction 的对象来获取有关 Metal 的函数库以及它所包含的函数的引用了。(这个概念我们在 Metal#1 Initialization 中也提及过)。
然后如果我们要获得一个关于某个函数的信息(比如说这个 add_arrays),就让 MTLDevice 创建一个 MTLLibrary 对象,然后叫这个对象返回一个指代该着色器函数的 MTLFunction 对象。
我:Device,能不能给我 add_arrays?
Device:我去问问Library。(转头)Library,给下add_arrays。
Library:你好,这是 add_arrays,给您打包了,装在 MTLFunction 袋子里了哈。
Device:哦好。(转头)来,这是 add_arrays。
我:okk,感谢。
这很好理解吧。把上面这段变成代码,就是获取 add_arrays 的方法了。
// 我:Device,能不能给我 add_arrays?
- (instanceType) initWithDevice: (id<MTLDevice>) device
{
self = [super.init];
if (self) {
{
_mDevice = device;
NSError *error = nil;
// Device:我去问问Library。
id<MTLLibrary> defaultLibrary = [_mDevice newDefaultLibrary];
// (转头)Library,给下add_arrays。
if (defaultLibrary == nil) {
NSLog(@"Failed to find the default lib.");
return nil;
}
// Library:你好,这是 add_arrays,给您打包了,装在 MTLFunction 袋子里了哈。
id<MTLFunction> addFunction = [defaultLibrary newFunctionWithName: @"add_arrays"];
if (addFunction == nil) {
NSLog(@"Failed to find the adder function.");
return nil;
}
// Device:哦好。(转头)来,这是 add_arrays。
// 我:okk,感谢。
}
渲染管线修改
在 Metal#2 Rendering Pipeline 中我们已经介绍过与 Metal 渲染管线相关的概念。在这里我们再复习一下以下概念之间的关系:管线、管线状态、指令队列、指令缓冲、指令编码、通道。
参考资料:
Comments