对thread_position_in_grid感到困惑答案

【问题标题】：Confused about thread_position_in_grid对thread_position_in_grid感到困惑
【发布时间】：2019-09-13 05:50:49
【问题描述】：

我正在 macOS 上的 Metal 中开发计算着色器。我正在尝试做一些非常基本的事情来了解它们是如何工作的。我看到一些我不明白的输出。我想我会先尝试生成一个简单的 2D 渐变。红色通道沿宽度从 0 增加到 1，绿色通道沿高度从 0 增加到 1。所以我写了这个内核：

kernel void myKernel(texture2d<half, access::write> outTexture [[ texture(MBKT_OutputTexture) ]],
                     uint2  gid  [[thread_position_in_grid]])
{
    half4  color = half4((float)gid.x / 480.0, (float)gid.y / 360.0, 0.0, 1.0);

    outTexture.write(color, gid);
}

我得到的是在中途点从 0 增加到 0.5，而图像的其余部分是稳定的 0.5，如下所示：

如果我反转 2 个值以便内核计算：

half4  color = half4(1.0 - (float)gid.x / 480.0, 1.0 - (float)gid.y / 360.0, 0.0, 1.0);

结果更奇怪。我希望它在左侧和底部是 1.0，在中间下降到 0.5，但相反，我得到了这个：

这里发生了什么？在第一种情况下，就像所有超过中点的值都为 0.5。在第二种情况下，就像左/下边缘为 0.5，中间为 1.0，然后在一个像素后翻转回 0.0。

奇怪的是，如果我使用 thread_position_in_grid 从缓冲区中提取值，它可以正常工作。例如，我可以计算一个 Mandelbrot 集，结果是正确的。但是我对上面的简单内核发生的事情感到困惑。谁能给我解释一下？

这是我在MTKViewDelegate 中的计算内核设置代码。这是基于 Apple 的“Hello Compute”示例代码：

    _metalView = metalView;
    _device = metalView.device;
    _commandQueue = [_device newCommandQueue];

    _metalView.colorPixelFormat = MTLPixelFormatBGRA8Unorm_sRGB;

    // Load all the shader files with a .metal file extension in the project
    id<MTLLibrary> defaultLibrary = [_device newDefaultLibrary];

    // Load the kernel function from the library
    id<MTLFunction> kernelFunction = [defaultLibrary newFunctionWithName:@"myKernel"];

    // Create a compute pipeline state
    NSError*    error   = nil;
    _computePipelineState = [_device newComputePipelineStateWithFunction:kernelFunction
                                                                   error:&error];

    if(!_computePipelineState)
    {
        NSLog(@"Failed to create compute pipeline state, error %@", error);
        return nil;
    }

这是我创建输出纹理和线程组的代码：

MTLTextureDescriptor*   outputTextureDescriptor = [MTLTextureDescriptor texture2DDescriptorWithPixelFormat:MTLPixelFormatBGRA8Unorm_sRGB
                                                                                                     width:_viewportSize.x
                                                                                                    height:_viewportSize.y
                                                                                                 mipmapped:NO];
_outputTexture = [_device newTextureWithDescriptor:outputTextureDescriptor];

// Set the compute kernel's threadgroup size of 16x16
_threadgroupSize = MTLSizeMake(16, 16, 1);

// Calculate the number of rows and columns of threadgroups given the width of the input image
// Ensure that you cover the entire image (or more) so you process every pixel
_threadgroupCount.width  = (_viewportSize.x + _threadgroupSize.width - 1) / _threadgroupSize.width;
_threadgroupCount.height = (_viewportSize.y + _threadgroupSize.height - 1) / _threadgroupSize.height;

// Since we're only dealing with a 2D data set, set depth to 1
_threadgroupCount.depth = 1;

在我的测试中，_viewportSize 是 480 x 360。

我在 cmets 中完成了 @Egor_Shkorov 建议的额外测试。我没有硬编码 480 和 360，而是使用了 threads_per_grid 变量：

kernel void myKernel(
                             texture2d<half, access::write> outTexture [[ texture(MBKT_OutputTexture) ]],
                             uint2  gid  [[thread_position_in_grid]],
                             uint2 tpg [[threads_per_grid]])
{

    half4  color = half4((float)gid.x / tpg.x, (float)gid.y / tpg.y, 0.0, 1.0);

    outTexture.write(color, gid);
}

这改进了事情，使梯度在每个方向上一直拉伸，但它仍然只从 0 到 0.5，而不是在每个方向上到 1：

【问题讨论】：

您需要显示调度此计算着色器的应用程序代码，尤其是线程组大小和线程（组）计数。另外，纹理是如何创建的。（并在您的计算着色器中显示outTexture 的声明。始终显示真实代码，因为编辑后的代码可能不代表您的实际问题。）
好的，我已经在上面添加了。如果有任何遗漏，请告诉我。
我建议使用threads_per_grid 而不是硬编码值，然后检查输出是否相同。
有趣！这会导致各种图块正确排列，因此我得到了从左到右和从上到下的平滑渐变，但不是在每个方向上都得到从 0 到 1 的渐变，它仍然是从 0 到 0.5 的渐变。为什么？
如果你这样做了half4 color = half4((float)gid.x / (float)outTexture.get_width(), (float)gid.y / (float)outTexture.get_height(), 0.0, 1.0);。此外，您应该检查gid.x 和gid.y 永远不会大于输出纹理的宽度/高度，否则您最终会在纹理内存之外写入，并且会发生不好的事情。（注意 360 不是 16 的整数倍。）

标签： metal compute-shader

【解决方案1】：

非常相似的事情发生在我身上。 thread_position_in_grid 的值似乎被限制在一个小范围内，而不是整个网格（可能只有threadgroup 的大小）。总之，可能是因为你在打电话

_commandEncoder.dispatchThreads(threadGroupCount, threadsPerThreadgroup: threadGroupSize)

而不是

_commandEncoder.dispatchThreadgroups(threadGroupCount, threadsPerThreadgroup: threadGroupSize)

我注意到属性thread_position_in_grid 在这些函数下会产生不同的值。不确定这是否是预期行为，因为我在文档中找不到相关描述，并且我希望该属性指的是整个网格中的位置。此外，Metal 在使用dispatchThreads() 时会决定线程组的数量，并且可以创建可能与问题有关的非统一线程组。

dispatchThreads(_:threadsPerThreadgroup:)

仅当设备支持非均匀时才使用此方法线程组大小。见Metal Feature Set Tables。该方法编码一个调度调用，指定网格中任意数量的线程 (threadsPerGrid)。 Metal 计算所需线程组的数量，如有必要，提供部分线程组。当计算命令已编码，对参数或资源的任何必要引用之前在编码器上设置的内容会被记录为命令的一部分。对命令进行编码后，您可以安全地将编码状态更改为设置编码其他命令所需的参数。

【讨论】：