金属着色语言 - 缓冲区绑定

Question

我想要粒子随时间增加。我得到将缓冲区值设置得更高，这样我就可以调整粒子的数量。我在想的是我将为缓冲区设置最大计数大小，然后在 shader 中，我将有一个带有数组的 struct 来获取粒子属性。

我的swift里有这个：

var vectMaxCount = 10
var metalvects = [float3(0.0,0.0,0.0),float3(1.0,0.0,0.0),float3(2.0,0.0,0.0)]
var vectBuffer: MTLBuffer!

然后我注册 buffer:

vectBuffer  = device!.makeBuffer(length: MemoryLayout<float3>.size * vectMaxCount, options: [])

并相应地更新 buffer：

...
command_encoder.setBuffer(vectBuffer, offset: 0, at: 2)
var bufferPointer = vectBuffer.contents()
memcpy(bufferPointer, &metalvects, MemoryLayout<float3>.size * vectMaxCount)

let threadGroupCount = MTLSizeMake(8, 8, 1)
let threadGroups = MTLSizeMake(drawable.texture.width / threadGroupCount.width, drawable.texture.height / threadGroupCount.height, 1)
command_encoder.dispatchThreadgroups(threadGroups, threadsPerThreadgroup: threadGroupCount)
command_encoder.endEncoding()
command_buffer.present(drawable)
command_buffer.commit()

并尝试从 metal 文件中获取它：

struct Vects
{
    float3 position[100];
};

kernel void compute(texture2d<float, access::write> output [[texture(0)]],
                    constant Vects &vects [[buffer(2)]],
                    uint2 gid [[thread_position_in_grid]]) {
...
}

我得到一个错误：

validateComputeFunctionArguments:727: failed assertion `(length - offset)(160) must be >= 1600 at buffer binding at index 2 for vects[0].'

提示command_encoder.dispatchThreadgroups(threadGroups, threadsPerThreadgroup: threadGroupCount)行报错。我阅读了一些关于 buffer binding 的内容，我认为是我发送 threadGroupCounts 或 ThreadGroup 的方式给我带来了问题。

如果我将 float3 position[100]; 更改为 float3 position[7];，它仍然有效。超过 7 的任何值都会出现类似的错误。

我该如何解决这个问题？

有没有好的公式来估计threadGroups和threadGroupCount？即使是经验法则也能做到这一点？

Update01

根据 Ken Thomases 的回答，我将代码更改为：

swift:

vectBuffer  = device!.makeBuffer(length: MemoryLayout<float3>.stride * metalvects.count, options: [])
...
memcpy(bufferPointer, &metalvects, MemoryLayout<float3>.stride * metalvects.count)
...

金属：

struct Vects
{
    float3 position[3];
};
...

目前确实有效。但是，我如何分配更高的缓冲内存，以便稍后在应用程序中使用，如提到的？

Answer 1

这里有多个问题。

您正在用特定大小定义 Vects。这允许 Metal 检查索引 2 处的缓冲区大小是否足够大以匹配 vects 变量的大小。它在抱怨，因为它不够大。（例如，如果 vects 被声明为 constant float3 *vects [[buffer(2)]]，它将无法执行此检查。）

其次，您的缓冲区大小 — MemoryLayout<float3>.size * vectMaxCount — 不正确。它没有考虑 float3 的对齐方式，因此没有考虑 [float3] 数组中元素之间存在的填充。正如 documentation for MemoryLayout 中所述，在计算分配大小时，您应该始终使用 stride，而不是 size。

这就是当 Vects::position 的长度为 8 个或更多元素时失败的原因。你会期望它从 11 个元素开始，因为 vectMaxCount 是 10，但是你的缓冲区比 vectMaxCount float3 的数组短。具体来说，你的缓冲区是 10 * 12 == 120 字节长。 float3 的步幅为 16 和 120 / 16 == 7.5.

如果您在分配缓冲区时从 size 切换到 stride 并将 Vects::position 的元素计数更改为 10 以匹配 vectMaxCount，那么您将得到过去这个紧迫的问题。但是，还有其他潜伏的问题。

您当前的计算函数不知道 vects.position 中实际填充了多少元素。您需要传入元素的实际数量。

这一行：

memcpy(bufferPointer, &metalvects, MemoryLayout<float3>.size * vectMaxCount)

不正确（即使在将 size 替换为 stride 之后）。它读取了 metalvects 的末尾。那是因为 metalvects 中的元素数量少于 vectMaxCount。您应该使用 metalvects.count 而不是 vectMaxCount.

金属着色语言 - 缓冲区绑定

Metal Shading Language - buffer binding

kernel

metal

macos-sierra