内存写入性能 – 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数据

对着色器的每次更新都会以结构中一对floattypes的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 。 这取决于你的数据。