top of page
  • Writer's pictureLingheng Tao

Metal #21 [Appendix] Compute Shaders



这篇笔记主要写一下关于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 渲染管线相关的概念。在这里我们再复习一下以下概念之间的关系:管线、管线状态、指令队列、指令缓冲、指令编码、通道。


 

参考资料:

18 views0 comments

Recent Posts

See All

Comments


bottom of page