ios - 在 Metal 中填充 float 缓冲区

问题:

我需要用一个常量值填充 Float 的 MTLBuffer，比如 1729.68921。我还需要它尽可能快。

因此，我被禁止在 CPU 端填充缓冲区(即从 MTLBuffer 获取 UnsafeMutablePointer 并以串行方式分配)。

我的方法

理想情况下，我会使用 MTLBlitCommandEncoder.fill() ，但是 AFAIK 它只能用 UInt8 值填充缓冲区(假设 UInt8 是 1 字节长， Float 是 4 字节长，我不能指定我的 Float 常量的任意值)。

到目前为止，我只看到了 2 个选项，但似乎都过分了:

1. 创建另一个缓冲区 B 填充常量值并通过 MTLBlitCommandEncoder
将其内容复制到我的缓冲区中
2. 创建一个 kernel 函数来填充缓冲区

问题

用 Float 填充 MTLBuffer 的最快方法是什么 恒定值？

---

Best Answer-推荐答案

使用从每个线程写入多个缓冲区元素的计算着色器是我实验中最快的方法。这取决于硬件，因此您应该在希望部署应用的所有设备上进行测试。

我编写了两个计算着色器:一个填充 16 个连续数组元素而不检查数组边界，另一个在检查缓冲区长度后设置单个数组元素:

kernel void fill_16_unchecked(device float *buffer  [[buffer(0)]],
                              constant float &value [[buffer(1)]],
                              uint index            [[thread_position_in_grid]])
{
    for (int i = 0; i < 16; ++i) {
        buffer[index * 16 + i] = value;
    }
}

kernel void single_fill_checked(device float *buffer         [[buffer(0)]],
                                constant float &value        [[buffer(1)]],
                                constant uint &buffer_length [[buffer(2)]],
                                uint index                   [[thread_position_in_grid]])
{
    if (index < buffer_length) {
        buffer[index] = value;
    }
}

如果您知道缓冲区计数始终是线程执行宽度乘以您在循环中设置的元素数量的倍数，则可以使用第一个函数。第二个函数是当你调度一个网格时的后备，否则会溢出缓冲区。

一旦您从这些函数构建了两个管道，您就可以使用一对计算命令来分派(dispatch)工作，如下所示:

NSInteger executionWidth = [unchecked16Pipeline threadExecutionWidth];
id<MTLComputeCommandEncoder> computeEncoder = [commandBuffer computeCommandEncoder];
[computeEncoder setBuffer:buffer offset:0 atIndex:0];
[computeEncoder setBytes:&value length:sizeof(float) atIndex:1];
if (bufferCount / (executionWidth * 16) != 0) {
    [computeEncoder setComputePipelineState:unchecked16Pipeline];
    [computeEncoder dispatchThreadgroups:MTLSizeMake(bufferCount / (executionWidth * 16), 1, 1)
                   threadsPerThreadgroup:MTLSizeMake(executionWidth, 1, 1)];
}
if (bufferCount % (executionWidth * 16) != 0) {
    int remainder = bufferCount % (executionWidth * 16);
    [computeEncoder setComputePipelineState:checkedSinglePipeline];
    [computeEncoder setBytes:&bufferCount length:sizeof(bufferCount) atIndex:2];
    [computeEncoder dispatchThreadgroups:MTLSizeMake((remainder / executionWidth) + 1, 1, 1)
                   threadsPerThreadgroup:MTLSizeMake(executionWidth, 1, 1)];
}
[computeEncoder endEncoding];

请注意，以这种方式完成工作并不一定比每个线程只写入一个元素的简单方法更快。在我的测试中，它在 A8 上快 40%，在 A10 上大致相当，在 A9 上慢 2-3 倍(!)。始终使用您自己的工作量进行测试。

关于ios - 在 Metal 中填充 float 缓冲区，我们在Stack Overflow上找到一个类似的问题： https://stackoverflow.com/questions/41198260/