新建计算管线状态与函数失败。

Question

我正在尝试在Metal上运行神经网络。

基本思想是数据复制。每个GPU线程为随机数据点运行一个版本的网络。

我已经编写了其他工作良好的着色器。

我还尝试在C++命令行应用程序中运行我的代码。那里没有错误。

我使用苹果文档将代码转换为Metal C ++，因为不支持来自C ++ 11的所有内容。

当它尝试将newComputePipelineStateWithFunction分配给金属设备时，它在加载内核函数后崩溃。这意味着存在问题的代码在编译时未被捕获。

MCVE：

kernel void net(const device float *inputsVector [[ buffer(0) ]], // layout of net *
                uint id [[ thread_position_in_grid ]]) {

    uint floatSize = sizeof(tempFloat);
    uint inputsVectorSize = sizeof(inputsVector) / floatSize;

    float newArray[inputsVectorSize];


    float test = inputsVector[id];

    newArray[id] = test;

}

更新

这与动态数组有很大关系。

由于无法创建管线状态并且在实际着色器运行时不崩溃，因此必须是编码问题，而不是输入问题。

从动态数组中分配值到缓冲区会导致失败。

- R Menke

---

你问题中的文字量令人惊恐。我确实读完了所有内容，但仍无法回答非常基本的问题：最小可重现代码示例在哪里？你所说的“不起作用”是指什么？ - Codeguard

你应该用最少的代码替换所有内容，以重现错误，并提供准确的错误描述。 - Codeguard

你熟悉Metal吗？那么你知道目前无法准确定位导致错误的代码行，或获取详细状态并找到破坏程序的值。不管怎样，我在解决这个问题一天后才问了这个问题。我不断添加我尝试过的和发现的内容。很抱歉文本量有点多，但由于这是全新的领域，添加更多信息会更好。 - R Menke

什么是悬赏奖励？我如何开始一个？如果您提出了一个好问题，并且通过状态和进展更新进行了编辑，但仍然没有得到答案，您可以通过在上面设置悬赏来引起关注。 - R Menke

从你的个人资料来看，我猜你不会为 Mac 编写代码。所以你不知道 Metal 可能会有多么痛苦。如果我们能保持话题不偏离主题，而不是让它变成形式问题，我会很感激的。如果有一天苹果公司能够让我们理解为什么某个管线失败了，那么就会有很多好的问题。但在那之前，这将是混乱、令人沮丧且没有太多答案的。也许我会很幸运，有人已经遇到过类似的问题。 - R Menke

Answer 1

真正的问题： 这是一个内存问题！

对于所有说这是内存问题的人，你们是正确的！以下是一些伪代码来说明这个问题。抱歉它是用“Swift”编写的，但易于阅读。Metal Shaders 有一种奇怪的方式来运行。它们首先被初始化而没有值以获取内存。失败的原因在于这一步依赖于后面的步骤：设置缓冲区。

所有问题都归结于哪些值何时可用。我对 newComputePipelineStateWithFunction 的理解是错误的。它不仅仅是获取着色器函数，它也是初始化过程中的一个微小步骤。

class MetalShader {

    // buffers
    var aBuffer : [Float]
    var aBufferCount : Int

    // step One : newComputePipelineStateWithFunction
    memory init() {
        // assign shader memory

        // create memory for one int
        let aStaticValue : Int
        // create memory for one int
        var aNotSoStaticValue : Int // this wil succeed, assigns memory for one int

        // create memory for 10 floats
        var aStaticArray : [Float] = [Float](count: aStaticValue, repeatedValue: y) // this will succeed

        // create memory for x floats
        var aDynamicArray : [Float] = [Float](count: aBuffer.count, repeatedValue: y) // this will fail
        var aDynamicArray : [Float] = [Float](count: aBufferCount, repeatedValue: y) // this will fail

        let tempValue : Float // one float from a loop

    }

    // step Two : commandEncoder.setBuffer()
    assign buffers (buffers) {

        aBuffer = cpuMemoryBuffer

    }

    // step Three : commandEncoder.endEncoding()
    actual init() {
        // set shader values

        let aStaticValue : Int = 0

        var aNotSoStaticValue : Int = aBuffer.count

        var aDynamicArray : [Float] = [Float](count: aBuffer.count, repeatedValue: 1) // this could work, but the app already crashed before getting to this point.

    }

    // step Four : commandBuffer.commit()
    func shaderFunction() {
        // do stuff
        for i in 0..<aBuffer.count {

            let tempValue = aBuffer[i]

        }
    }
}

修复：

我终于意识到缓冲区在技术上是动态数组，而不必在着色器内创建数组，我也可以添加更多的缓冲区。这显然有效。

- R Menke

Answer 2

我认为你的问题出在这一行：

uint schemeVectorSize = sizeof(schemeVector) / uintSize;

这里的schemeVector是动态的，所以就像经典的C++一样，你不能使用sizeof来获取动态数组的元素数量。 sizeof只能在你在金属着色器代码中定义本地/静态数组时才能工作。

想象一下它在内部是如何工作的：在编译时，Metal编译器应该将sizeof调用转换为常量...但由于schemeVector是您着色器的参数，因此可以具有任何大小，所以它无法这样做...

因此，对我来说解决方案是在您的C++/ObjectiveC/Swift代码中计算schemeVectorSize，并将其作为参数传递给着色器（在OpenGLES术语中作为统一变量...）。

- VB_overflow

---

我对编译器的工作原理以及苹果是如何决定使其不受许多易于捕获的错误影响进行了一些研究。我认为代码之所以不被视为错误，是因为你可以静态声明一个uint并使用它来设置多个数组的大小。但我认为newComputePipelineStateWithFunction也会为着色器中声明的所有内容保留内存。因此将schemeVectorSize作为参数传递不起作用。它仍然不知道为代码的非缓冲区部分保留多少内存。 - R Menke

我想说的是 uint schemeVectorSize = sizeof(schemeVector) / uintSize; 等同于 uint schemeVectorSize = sizeof(const device uint*) / uintSize;，这可能会导致 schemeVectorSize 等于1或2...所以这可能完全不是你期望的结果...你尝试使用离线Metal编译器编译着色器以查看是否有错误了吗？ - VB_overflow

正确的答案应该是“可以”。您可以获取缓冲区的大小并将其用于循环或if语句。您不能使用它来构建动态数组。因为据我所知，除了缓冲区之外的着色器中的每个实例都必须在分配缓冲区之前创建。 - R Menke

我总是使用离线编译器。那里没有错误。请看我的第一条评论，以了解原因。 - R Menke

Menke，你不能在着色器代码内部调用sizeof(pointer)并期望它正常工作。你必须将常量缓冲区大小作为常量空间值传递到着色器调用中，通常包装在一个结构体中。着色器代码不知道稍后传递到着色器中的缓冲区的长度。 - MoDJ