iOS Metal Performance Shaders 图像上采样答案

【问题标题】：iOS Metal Performance Shaders Image UpsampleiOS Metal Performance Shaders 图像上采样
【发布时间】：2017-01-25 08:35:15
【问题描述】：

对于带有 Metal Performance Shader 的 iOS 应用，我想为 CNN 平均池化层反向传播编写一个 GPU 加速函数。这与图像上采样几乎相同。例如，如果输入是

2 5 6
3 6 7
8 9 0

上采样图像应该是

2 2 5 5 6 6
2 2 5 5 6 6
3 3 6 6 7 7
3 3 6 6 7 7
8 8 9 9 0 0
8 8 9 9 0 0

我编写了以下内核函数：

kernel void upsample(texture2d<float, access::read> inTexture [[texture(0)]],
                     texture2d<float, access::write> outTexture [[texture(1)]],
                     uint2 gid [[thread_position_in_grid]])
{
    const float4 color = inTexture.read(gid);

    uint2 p;
    p = uint2(gid.x * 2, gid.y * 2);
    outtexture.write(color, p);
    p = uint2(gid.x * 2 + 1, gid.y * 2); 
    outtexture.write(color, p);
    p = uint2(gid.x * 2, gid.y * 2 + 1);
    outtexture.write(color, p);
    p = uint2(gid.x * 2 + 1, gid.y * 2 + 1);  
    outtexture.write(color, p);
}

但我不确定这是否正确。我怎么知道原始的“gid”与输入纹理中的坐标有关，但与输出纹理中的坐标无关？

【问题讨论】：

附带说明，如果您希望通过大型数据集最大限度地提高性能，我的经验表明，您可以通过使用固定 GPU 硬件来实现显着提高的性能。只需设置渲染管道，加载纹理，配置适当大小的渲染目标，将纹理采样器设置为“最近”放大率，然后让固定硬件进行缩放。在我进行的其他类似测试中，我发现固定硬件做这种事情的速度比计算内核快几倍。

标签： ios gpu metal

【解决方案1】：

这种需求现在由 MPSCNNUpsampling 系列内核处理。请参阅 MetalPerformanceShaders/MPSCNNUpsampling.h

【讨论】：

我知道这是在 8 月，但这并不一定能完成工作，就好像 OP 只想使用 MTLTexture 一样，仅为上采样进行转换可能是一项昂贵的操作。

【解决方案2】：

@消防员我对您的上采样代码进行了概要分析，但是一一写入 4 个数据的成本很高。还有你关于gid的问题。它与输入和输出纹理的内核进程区的工作项有关。也许它在输入中包含 1 个像素，而在输出中包含 4 个像素。

【讨论】：

【解决方案3】：

假设您发送的网格与源纹理的尺寸相同，那么您所写的内容对我来说是正确的。

网格中的线程位置与您读取或写入的坐标之间没有固有的对应关系。 gid 只是您使用计算命令编码器分派的工作项网格中的一个点。在您的内核中，您可以随心所欲地解释它。如果您调度一个 2D 网格并且您正在处理纹理，那么将gid 视为一对坐标通常是有意义的，就像您在此处所做的那样。

您也可以自己调度 1D 网格并将其“包裹”到 2D 网格，但使用 2D 线程位置可以在内核中节省（昂贵的）整数除法，而且更直观。

【讨论】：