内存写入性能 – GPU CPU共享内存
我将根据MTLBuffer
提供的共享GPU / CPU文档,使用posix_memalign
来分配input和输出MTLBuffer。
另外:使用最新的API比posix_memalign
更容易
let metalBuffer = self.metalDevice.newBufferWithLength(byteCount, options: .StorageModeShared)
我的内核函数在大约一千六百万个复杂的价值结构上运行,并写出相同数量的复数值结构来存储。
我已经执行了一些实验,我的金属内核'复杂的math部分'在0.003秒内执行(是!),但是将结果写入缓冲区需要> 0.05(No!)秒。 在我的实验中,我注释了math部分,只是把零分配给内存,它需要0.05秒,注释掉这个任务,并把算术加回去0.003秒。
在这种情况下共享内存是否缓慢,或者我可能会尝试一些其他的提示或技巧?
额外的细节
testing平台
- iPhone 6S – 每帧约0.039秒
- iPad Air 2 – 每帧〜0.130秒
stream数据
对着色器的每次更新都会以结构中一对float
types的forms接收大约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个通道×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上。 似乎瓶颈是写入input缓冲区的值。 那么,我想这个问题就成了如何避免瓶颈?
内存拷贝
我写了一个快速testing来比较各种字节复制方法的性能。
当前状态
我已经将执行时间减less到了0.02秒,这听起来不是很多,但是每秒的帧数却有很大的不同。 目前最大的改进是切换到cblas_scopy()
。
减lesstypes的大小
最初,我是预先转换签名的16位整数作为浮动(32位),因为最终这是如何使用它们。 这是性能开始迫使您将值存储为16位以将您的数据大小减半的情况。
Objective-C超过Swift
对于处理数据移动的代码,你可以selectObjective-C over Swift(Warren Moore推荐)。 Swift在这些特殊情况下的性能仍然没有达到要求。 您也可以尝试调用memcpy
或类似的方法。 我见过一些使用循环缓冲指针的例子,这个在我的实验中慢慢地进行了。
testing困难
我真的想在机器的操场上进行一些关于各种复制方法的实验,不幸的是这种做法没用。 相同实验的iOS设备版本执行完全不同。 有人可能会认为相对的performance会是相似的,但是我觉得这也是一个无效的假设。 如果你有一个使用iOS设备作为解释器的游乐场,那将非常方便。
通过将数据编码为huffman代码并在GPU上进行解码,您可能会获得大幅度的加速,请参阅MetalHuffman 。 这取决于你的数据。