[MetalKit]46-Introduction to compute using Metal 用 Metal 进行计算的简介

本系列文章是对 metalkit.org 上面MetalKit内容的全面翻译和学习.

MetalKit系统文章目录

在 GPU 编程的领域中，计算或者说GPGPU，是 GPU 编程中除渲染外的另一种用途。它们都涉及到了 GPU 并行编程，不同之处在于在计算中对线程的工作方式进行了更精细的控制。这样，当你想要某些线程来处理问题的某一部分，同时其他线程去处理该问题的另一部分时，就会很有用。

本文是一系列关于计算的文章的开始篇。本文中的主题是关于图像处理，因为它是引入计算和线程管理的最简单方法。

注意：本文假设您知道如何创建一个微型的Metal项目或playground，可以将屏幕清除为纯色。

第一个不同点就是，你需要创建一个MTLComputePipelineState以取代以前渲染时用的MTLRenderPipelineState：

let function = library.makeFunction(name: "compute")
let pipelineState = device.makeComputePipelineState(function: function)

第二件事是，你需要一个纹理，以供线程使用。如果你使用的是playground，那你只需要下面几行：

let textureLoader = MTKTextureLoader(device: device)
let url = Bundle.main.url(forResource: "nature", withExtension: "jpg")!
let image = try textureLoader.newTexture(URL: url, options: [:])

第三件事，你需要一个MTLComputeCommandEncoder对象，以便将先前创建的管线状态对象和纹理，都附着上去：

commandEncoder.setComputePipelineState(pipelineState)
commandEncoder.setTexture(image, index: 0)

第四件事，你需要一个kernel shader内核着色器，要记得，你之前开始时就为其创建了一个名为compute的函数。当然，你可以将内核代码放到 .metal文件里：

kernel void compute(texture2d<float, access::read> input [[texture(0)]],
                    texture2d<float, access::write> output [[texture(1)]],
                    uint2 id [[thread_position_in_grid]]) {
    float4 color = input.read(id);
    output.write(color, id);
}

在着色代码中，输入是你先前创建的MTLTexture对象，称为image，输出是一个可绘制纹理，你将向其中写入数据，然后就可以被呈现到屏幕上了：

let drawable = view.currentDrawable
commandEncoder.setTexture(drawable.texture, index: 1)

第五件事也是最后一件事是，你需要调度线程来干活。有趣的事情就从现在开始了！你需要做的是在commandEncoder中结束编码之前，加上几句代码：

let threadsPerGroup = MTLSizeMake(100, 10, 1)
let groupsPerGrid = MTLSizeMake(15, 90, 1)
commandEncoder.dispatchThreadgroups(groupsPerGrid, threadsPerThreadgroup: threadsPerGroup)

那么这里是怎么做的呢？线程是以网格(grid)形式来调度处理数据的，网格可以是 1-，2-，或3-维的。在本例中，你用的是 2D 的网格，因为要处理的是一张图片。不考虑维度的话，网格总是分割成多个线程组的，如下面的公式：

gridSize = groupsPerGrid * threadsPerGroup

在本例中，你定义一个组有100 x10个线程，每个网格有15 x 90组。如果你运行你的 playground，你会看到类似下面的情况：

边上的红色是什么东西？这是因为你试图去猜测图片的尺寸大小而导致的问题，线程数和组数应该用更“聪明”的方式获取。

显然，图像在两个维度上都大于分派的线程数。您可以做的一件事是使用图像大小进行有根据的猜测，以获得真正应该使用的组数量：

let width = Int(view.drawableSize.width)
let height = Int(view.drawableSize.height)
let w = threadsPerGroup.width
let h = threadsPerGroup.height
let groupsPerGrid = MTLSizeMake(width / w, height / h, 1)

运行一下，图片看起来会好很多了：

这里又出现一个新的问题---利用不足。请看下图的图表：

通常，您会认为正确设计的网格是3 x 2组，每组4 x 4个线程，因此网格为12 x 8个线程。然而，底部和右侧边缘的一些螺纹未得到充分利用，因为它们没有工作要做。

如果你制作一个较小的网格，比如8 x 4，它将会填满整个组，又会产生你在开始时看到的红色条带。这意味着唯一可接受的解决方案是修复未充分利用问题。您可以通过在每个维度中添加额外的组来解决此问题，如下所示：

let groupsPerGrid = MTLSizeMake((width + w - 1) / w, (height + h - 1) / h, 1)

你所做的就是用(w-1, h-1, 1)来实际扩大网格尺寸。这又带来了另一个风险 --- 访问越界坐标。要处理这个问题，您需要在读取输入图像之前向内核着色器添加边界检查：

if (id.x >= output.get_width() || id.y >= output.get_height()) {
    return;
}

这将处理那些不应该做任何工作的线程，并处理越界的访问。

那个线程组的大小怎么样 --- 无法优化吗？到目前为止，你一直在猜这些尺寸。当然，还有一种方法可以获得最佳的群组尺寸。硬件提供了一些可以通过管道状态对象(pipeline state object)访问的功能：

var w = pipelineState.threadExecutionWidth
var h = pipelineState.maxTotalThreadsPerThreadgroup / w
let threadsPerGroup = MTLSizeMake(w, h, 1)

线程执行宽度（在其他API中也称为wavefront或warp）是GPU组合在一起的线程数，因此它们可以并行地在不同的数据上执行相同的指令。组中的线程数应该是threadExecutionWidth的倍数，但绝不能大于maxTotalThreadsPerThreadgroup。

那太棒了！如何找到办法，来避免做这些未充分利用和边界检查呢？Metal 也在这里给你提供了帮助。无需使用dispatchThreadgroups()，API提供了更新的dispatchThreads()函数，它实现了两件事：

通过自动创建非均匀线程组（例如3 x 4）来适应边缘情况，这样就避免让你处理未充分利用的问题。
它甚至可以决定需要多少组，前提是您为其提供网格大小和您想要使用的组大小。

注意：dispatchThreads()函数适用于所有macOS设备，但它不适用于使用A10或更旧处理器的iOS设备。

你需要做的就是，就下面代码替换计算每个网格组数的代码：

w = Int(view.drawableSize.width)
h = Int(view.drawableSize.height)
let threadsPerGrid = MTLSizeMake(w, h, 1)
commandEncoder.dispatchThreads(threadsPerGrid, threadsPerThreadgroup: threadsPerGroup)

但是等一下，我是不是说过：这里是最好玩的地方？是的，然后来到 kernel shader 中，移除边界检查代码，因为现在已经不需要它了。然后在最后一行前，添加下面代码，倒转颜色通道：

color = float4(color.g, color.b, color.r, 1.0);

运行一下 playground，你会看到类似下面的图像：

将上一行用下面代码替换，它将灰度应用于图像：

color.xyz = (color.r * 0.3 + color.g * 0.6 + color.b * 0.1) * 1.5;

运行一下 playground，你会看到类似下面的图像：

最后，将下面代码替换：

float4 color = input.read(id);
color.xyz = (color.r * 0.3 + color.g * 0.6 + color.b * 0.1) * 1.5;

替换为下面的代码，这里将图片将图像像素化为5像素的正方形：

uint2 index = uint2((id.x / 5) * 5, (id.y / 5) * 5);
float4 color = input.read(index);

运行一下 playground，你会看到类似下面的图像：

玩得开心么？希望你玩得开心。如果你想要学习更多关于图像处理的知识，Simon Gladman有一本好书，Core Image For Swift。本文只是一个对 GPGPU 和GPU计算功能的简短介绍。请继续关注新主题。

源代码已经发布在Github上。本文基于书籍Metal by Tutorials的第 16 章完成。

下次见！