访问金属计算内核中网格位置以外的像素？答案

【问题标题】：Access pixels beyond grid position in metal compute kernel?访问金属计算内核中网格位置以外的像素？
【发布时间】：2019-12-16 13:37:56
【问题描述】：

我有一个金属核函数。通常你会像这样访问像素：

kernel void edgeDetect(texture2d<half, access::sample> inTexture [[ texture(0) ]],
                    texture2d<half, access::write> outTexture [[ texture(1) ]],
                    device const uint *roi [[ buffer(0) ]],
                    uint2 grid [[ thread_position_in_grid ]]) {

  if (grid.x >= outTexture.get_width() || grid.y >= outTexture.get_height()) {
      return;
  }

  half c[9];
  for (int i=0; i < 3; ++i) {
    for (int j=0; j < 3; ++j) {
      c[3*i+j] = inTexture.read(grid + uint2(i-1,j-1)).x;
    }
  }

  half3 Lx = 2.0*(c[7]-c[1]) + c[6] + c[8] - c[2] - c[0];
  half3 Ly = 2.0*(c[3]-c[5]) + c[6] + c[0] - c[2] - c[8];
  half3 G = sqrt(Lx*Lx+Ly*Ly);

  outTexture.write(half4(G, 0.0), grid);
}

现在我需要像这样访问当前网格位置附近的像素：

      half4 inColor = inTexture.read(grid - uint2(-1,-1));

基本上这是可行的，但是在线程边界上我有“不连续”，如图所示（砖墙图案）。

这很清楚，因为每个线程只传递它要处理的子纹理。所以超出线程边界我无法访问像素。

我的问题是：当我需要在计算内核中处理超出当前位置的像素时，这个概念是什么？这完全可以通过计算内核实现吗？

【问题讨论】：

我真的不明白你的意思是“这很清楚，因为每个线程都只传递它的子纹理来处理。”您能否充实不起作用的示例代码并提供一些入队参数？鉴于您发布的代码，我完全不清楚您将如何获得提供的图像 - 似乎缺少很多步骤，我不相信问题出在您认为的位置。
我已经添加了实际运行的代码！它在 (32,16,1) 的线程组中线程化，每个线程组有适当数量的线程。

标签： ios shader metal compute-shader

【解决方案1】：

我找到了问题：

线

c[3*i+j] = inTexture.read(grid + uint2(i-1,j-1)).x;

必须改为：

c[3*i+j] = inTexture.read(grid + uint2(i,j)).x;

显然，纹理中的 -1 位置索引失败，并产生了如上图所示的砖墙。

【讨论】：

【解决方案2】：

为了确保有人将其附加到此评论中作为答案：对于您可以在计算着色器中访问哪些像素没有任何限制。您的网格大小仅影响调度。

您的错误是将 unsigned uint2 实例化为 negative 数字。在循环的第一次迭代中，您将尝试构造 uint2(-1, -1)，它与 uint2(4294967295, 4294967295) 相同，因此方式超出范围。

您可以使用int2，或者根据您的自我回答，避免使用负数。

【讨论】：