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

Question

我分配输入和输出MTLBuffer使用posix_memalign根据共享GPU/CPU文档由memkite提供.

除此之外:使用最新的API比使用muck更容易 posix_memalign

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

我的内核函数在大约1600万个复杂值结构上运行,并向内存写出相同数量的复数值结构.

我已经完成了一些实验,我的Metal内核'复杂数学部分'在0.003秒内执行(是!),但是将结果写入缓冲区需要> 0.05(No!)秒.在我的实验中,我注释掉了数学部分并将零分配给内存,它需要0.05秒,注释掉分配并添加数学,0.003秒.

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

其他细节

测试平台

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

流数据

对着色器的每次更新都会在结构中以一对float类型的形式接收大约50000个复数.

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个通道x 130000个像素).我对每个'像素'执行一些操作,并对通道中的复杂结果求和,最后取复数的绝对值并将结果赋值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上.似乎瓶颈是将值写入输入缓冲区.那么,我想问题就是如何避免瓶颈？

记忆副本

我写了一个快速测试来比较各种字节复制方法的性能.

当前状态

我已经将执行时间减少到0.02秒,这听起来不是很多,但它在每秒帧数方面有很大差异.目前最大的改进是切换到的结果cblas_scopy().

Answer 1

减小字体大小

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

Objective-C 优于 Swift

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

测试难度

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