如何调试 Metal 计算着色器中难以理解的“内部错误”？（任意代码更改触发）

Question

我有一个适用于 iOS 的 Metal 计算着色器，它已开始生成： "Error Domain=AGXMetal Code=1 "Compiler encountered an internal error" newComputePipelineStateWithFunction() 期间出现错误。

这些错误在每次运行时都是一致的，但似乎是由对不相关代码的几乎任意修改触发的。我的意思是 - 在自然地尝试调试为什么我添加最新的代码块表面上导致了这个问题时，我发现删除看似任意且不相关的代码或构造行将消除错误。

我想知道我是否可能达到了编译器的某些大小或复杂性限制。

我的着色器函数总共不到 200 行代码，结构为几个 C 函数，并且不分配太多内存，但它确实有一些循环并传递一些缓冲区指针。在某种程度上，一切都工作得很好，最近添加的代码也更加相似。

我的问题是：

1）首先 - 计算管道的编译器到底在做什么（生成我的default.metallib时没有完成）以及是否有希望从中收集更多调试信息？

2）如果这是代码的某种大小或复杂性问题，是否有人知道我如何重组以减轻这种情况？这些有道理吗？

为此发布示例代码将很困难，但如果解决方案没有首先出现，我将尝试使用它进行更新。

编辑：

因此，我所做的就是煞费苦心地减少和简化我的代码，直到我有一个相对紧凑的示例来说明问题。这并不像听起来那么简单，因为许多看似微小的变化会导致问题消失，但当复杂性增加时，问题总是会回来。

请记住，这与下面的代码的作用无关在计算管道运行之前很长时间内设置计算管道时会发生故障。如果您发现明显错误的内容，请告诉我，但除此之外，该代码仅具有代表性。

下面的着色器在 A9 处理器（iPhone 6s 或 6s plus）上失败，但在 A7（iPad Air 第一代）上运行。

void myFunc( device int *ibuff0, thread int *ibuff1)
{
    int counter = 0;
    float fbuff0[8];
    for( int i = 0; i<8; i++) {
        if ( ibuff0[0] == 42 ) {
            fbuff0[counter++] = 0.0;
        }
    }

    float val = fbuff0[0];
    if ( distance( float2(0.0f,0.0f), float2(val,val) ) < 42.0f) {
        ibuff1[0] = 0;
    }
}

kernel void myKernelFunc( device int *ibuff0 [[ buffer(0) ]] )
{
    int ibuff1[8];
    myFunc( ibuff0, ibuff1 );
}

有趣的是有多少种方法可以解决上述问题。仅举几例：1）内联 myFunc（手动或使用 inline 关键字）。2) 注释掉任一缓冲区分配。3) 用本地缓冲区替换设备缓冲区。4）注释掉for循环，留下循环体。此外，距离函数调用在这里并不神奇，您可以替换那里使用“val”的任何非内联函数。

顺便说一句，这是一个完全伪造的单行版本，在 A9 和 A7 处理器上均失败：

kernel void myKernelFunc() {
    while ( true ) { }
}

还有一些想法——

我假设我在这里遇到的问题是 Metal 编译器尝试将这些条件和循环结构映射到可以在 GPU 上运行的动态统一代码类型时涉及的一些错误或限制。但我不知道为什么在撞到这堵墙之前我在代码中做到了这一点，因为上面的内容似乎并不比我成功做的事情更复杂。

现在我有了样本，我可以向 Apple 提交错误（正如一些人建议的那样）。但我想在这里分享，以防有人有想法。

更新：

我发现解决此问题的最简单方法是手动内联一些函数。