在金属中填充浮动缓冲区

Question

问题：

我需要补MTLBuffer的Float具有恒定值s -比方说1729.68921。我还需要尽快。

因此，禁止我填充CPU端的缓冲区（即UnsafeMutablePointer<Float>从MTLBuffer并以串行方式分配）。

我的方法

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

到目前为止，我只能看到2个选项，但这两个选项似乎都太过分了：

1. 创建另一个B填充有常量值的缓冲区，然后通过以下方式将其内容复制到我的缓冲区中MTLBlitCommandEncoder
2. 创建一个kernel将填充缓冲区的函数

问题

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

Answer 1

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

我编写了两个计算着色器：一个在不检查数组边界的情况下填充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;
    }
}

如果您知道缓冲区计数将始终是线程执行宽度乘以您在循环中设置的元素数的倍数，则可以使用第一个函数。第二个功能是当您分派可能会超出缓冲区范围的网格时的后备功能。

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

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倍（！）。始终根据自己的工作量进行测试。