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

Question

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

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

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

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

我已经进行了一些实验，我的Metal内核'复杂数学部分'在0.003秒内执行（是！），但是将结果写入缓冲区需要> 0.05（No！）秒。在我的实验中，我注释掉了数学部分并将零分配给内存，它需要0.05秒，注释掉分配并添加数学，0.003秒。

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

其他细节

Test platforms

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

The streaming data

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

struct ComplexNumber {
    float real;
    float imaginary;
};

Kernel signature

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。

Dispatch

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位大小的整数预转换为Floats（32位），因为最终它们将被如何使用。在这种情况下，性能开始强制您将值存储为16位以将数据大小减半。

Objective-C over Swift

对于处理数据移动的代码，您可以选择Objective-C over Swift（Warren Moore推荐）。 Swift在这些特殊情况下的性能仍然不尽如人意。您也可以尝试呼叫memcpy或类似的方法。我已经看过几个使用for-loop Buffer Pointers的例子，这在我的实验中表现得很慢。

测试困难

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

Answer 2

您可以通过将数据编码为霍夫曼代码并在GPU上解码来获得大幅加速，请参阅MetalHuffman。这取决于你的数据。