iOS金属性能着色器图像样本

Question

对于带有Metal Performance Shader的iOS App，我想为CNN平均池化层反向传播编写一个GPU加速函数。这和image upsample差不多。例如，如果输入是

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"与输入纹理中的坐标有关，而不是输出纹理?

Answer 1

你写的看起来对我来说是正确的，假设你分配了一个与源纹理具有相同尺寸的网格。

在网格中的线程位置和您读取或写入的坐标之间没有内在的对应关系。gid只是您使用计算命令编码器分派的工作项网格中的一个点。在内核中，您可以按照自己的意愿对其进行解释。如果你调度一个2D网格，你正在处理纹理，通常把gid当作一对坐标是有意义的，就像你在这里做的那样。

您也可以分派一个1D网格并将其"包装"到一个2D网格中，但是使用2D线程位置可以在内核中节省(昂贵的)整数除法，而且更直观。

Answer 2

@消防员我分析了你的抽样代码，但是一个接一个地写4个数据的成本很高。你关于gid的问题。它与输入和输出纹理的内核过程域的工作项相关。也许它包含1个输入像素和4个输出像素。

Answer 3

这个需求现在由MPSCNNUpsampling内核家族来处理。看到MetalPerformanceShaders/MPSCNNUpsampling.h