ios - 内存写入性能 - GPU CPU 共享内存

Question

我正在分配输入和输出 MTLBuffer使用 posix_memalign根据shared GPU/CPU documentation由 memkite 提供。

旁白:只使用最新的 API 比使用 posix_memalign 更容易

let metalBuffer = self.metalDevice.newBufferWithLength(byteCount, options: .StorageModeShared)

我的内核函数对大约 1600 万个复值结构进行运算，并将相同数量的复值结构写入内存。

我进行了一些实验，我的 Metal 内核“复杂数学部分”在 0.003 秒内执行(是!)，但将结果写入缓冲区需要 >0.05(否!)秒。在我的实验中，我注释掉了数学部分并将零分配给内存，这需要 0.05 秒，注释掉分配并添加数学部分，0.003 秒。

在这种情况下共享内存是否很慢，或者我是否可以尝试其他一些提示或技巧？

其他详细信息

测试平台

• iPhone 6S - 每帧约 0.039 秒
• iPad Air 2 - 每帧约 0.130 秒

流数据

着色器的每次更新都会收到大约 50000 个复数，形式为一对 float。在结构中键入。

struct ComplexNumber {
    float real;
    float imaginary;
};

内核签名

kernel void processChannelData(const device Parameters *parameters [[ buffer(0) ]],
                               const device ComplexNumber *inputSampleData [[ buffer(1) ]],
                               const device ComplexNumber *partAs [[ buffer(2) ]],
                               const device float *partBs [[ buffer(3) ]],
                               const device int *lookups [[ buffer(4) ]],
                               device float *outputImageData [[ buffer(5) ]],
                               uint threadIdentifier [[ thread_position_in_grid ]]);

除inputSampleData 外，所有缓冲区目前都包含不变的数据它接收我将要操作的 50000 个样本。其他缓冲区每个包含大约 1600 万个值(128 channel x 130000 像素)。我对每个“像素”执行一些操作并跨 channel 对复数结果求和，最后取复数的绝对值并分配结果 float至 outputImageData .

调度

commandEncoder.setComputePipelineState(pipelineState)

commandEncoder.setBuffer(parametersMetalBuffer, offset: 0, atIndex: 0)
commandEncoder.setBuffer(inputSampleDataMetalBuffer, offset: 0, atIndex: 1)
commandEncoder.setBuffer(partAsMetalBuffer, offset: 0, atIndex: 2)
commandEncoder.setBuffer(partBsMetalBuffer, offset: 0, atIndex: 3)
commandEncoder.setBuffer(lookupsMetalBuffer, offset: 0, atIndex: 4)
commandEncoder.setBuffer(outputImageDataMetalBuffer, offset: 0, atIndex: 5)

let threadExecutionWidth = pipelineState.threadExecutionWidth
let threadsPerThreadgroup = MTLSize(width: threadExecutionWidth, height: 1, depth: 1)
let threadGroups = MTLSize(width: self.numberOfPixels / threadsPerThreadgroup.width, height: 1, depth:1)

commandEncoder.dispatchThreadgroups(threadGroups, threadsPerThreadgroup: threadsPerThreadgroup)
commandEncoder.endEncoding()
metalCommandBuffer.commit()
metalCommandBuffer.waitUntilCompleted()

GitHub 示例

我写了一个例子叫做 Slow并把它放在 GitHub 上。似乎瓶颈是将值写入输入缓冲区。所以，我想问题变成了如何避免瓶颈？

内存复制

我写了一个quick test比较各种字节复制方法的性能。

现状

我已将执行时间减少到 0.02 秒左右，这听起来并不多，但它对每秒的帧数有很大影响。目前最大的改进是切换到 cblas_scopy() 的结果.

Answer 1

减少类型的大小

最初，我将带符号的 16 位大小的整数预先转换为 float (32 位)，因为最终这就是它们的使用方式。在这种情况下，性能开始迫使您将值存储为 16 位以将数据大小减半。

基于 Swift 的 Objective-C

对于处理数据移动的代码，您可以选择 Objective-C 而不是 Swift(Warren Moore 推荐)。 Swift 在这些特殊情况下的性能仍然达不到标准。您也可以尝试调用 memcpy 或类似的方法。我见过几个使用 for-loop 缓冲区指针的示例，但在我的实验中这执行起来很慢。

测试难度

我真的很想在机器上的 Playground 上做一些与各种复制方法相关的实验，不幸的是这没用。相同实验的 iOS 设备版本表现完全不同。有人可能认为相对性能会相似，但我发现这也是一个无效的假设。如果你能有一个使用 iOS 设备作为解释器的 Playground ，那将非常方便。