如何正确访问金属核函数中的邻居单元答案

【问题标题】：how to access the neighbor unit in metal kernel function correctly如何正确访问金属核函数中的邻居单元
【发布时间】：2016-07-22 08:48:45
【问题描述】：

我用金属做一些插值任务。我写的内核函数如下：

kernel void kf_interpolation( device short *dst, device uchar *src, uint id [[ thread_position_in_grid ]] )    
{
  dst[id] = src[id-1] + src[id] + src[id+1];
}

那个核函数不能给出预期值。我发现原因是 src[id-1] 始终为 0，这是错误值。但是， src[id+1] 包含正确的值。问题是我怎样才能正确使用邻居单元，例如[id-1]，在内核函数中。提前致谢。

【问题讨论】：

标签： ios macos shader metal

【解决方案1】：

处理这种边缘情况的最有效方法通常是在每一端增加源数组并偏移索引。因此，对于 N 次计算，请为您的 src 数组分配 N+2 个元素，用源数据填充元素 1 到 N（包括），并将元素 0 和 N+1 设置为您想要的边缘条件。

【讨论】：

【解决方案2】：

更有效的方法是使用 MTLTextures 而不是 MTLBuffers。 MTLTextures 有一个附加的寻址模式，当你读取纹理的边缘时，它会导致硬件自动替换零或最接近的有效纹素。他们还可以免费在硬件中进行线性插值，这对重采样有很大帮助，假设双线性插值对你来说已经足够好了。如果没有，我建议将 MPSImageLanczosScale 作为替代方案。

您可以从 MTLBuffer 制作 MTLTexture。两者将别名相同的像素数据。

【讨论】：

谢谢@Ollmann。 MTLTexture 是否适合存储视频编码中常见的 yuv 格式数据。
它可以做yuva 444没有太多麻烦。只需将它们粘贴在那里代替 RGBA。对于 422，MTLPixelFormatGBGR422 / MTLPixelFormatBGRG422 形式的支持有限。无论哪种情况，您都必须知道 RGB 通道中的数据实际上是 YCbCr，并在必要时在软件中执行您自己的 YCbCr ->RGB 转换矩阵 + 需要进行任何视频范围调整。
一种常见的 YUV 格式是 yuv420，它有 3 个通道：y、u 和 v。我可以为每个通道创建 3 个 MTLTexuture 对象，分别做计算任务吗？
如何实现线性“免费插值”？是用“采样器”完成的吗？你能在金属文档中给 for 一个关键字吗！谢谢
将过滤模式设置为线性而不是最近。如果坐标不是整数 + 0.5，则结果将是四个附近像素值的双线性插值。