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

Posted 2023-02-15

【中文标题】内存写入性能 - GPU CPU 共享内存

【英文标题】：Memory write performance - GPU CPU Shared Memory

【发布时间】：2016-06-22 00:23:08

【问题描述】：

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

另外：使用最新的 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 秒

流数据

对着色器的每次更新都会以结构中的一对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 上。似乎瓶颈是将值写入输入缓冲区。所以，我想问题变成了如何避免瓶颈？

内存拷贝

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

当前状态

我已将执行时间减少到 0.02 秒，这听起来并不多，但每秒的帧数却有很大的不同。目前最大的改进是切换到cblas_scopy()。

【问题讨论】：

您能以 GB/s 为单位描述您的内存带宽吗？您每帧写入多少字节并不明显。根据经验数据，我预计 iPhone 6 的一个微不足道的内核可以写入 0.5GB/s 到 1.5GB/s 之间，而 iPhone 6s 的写入速度大约是 iPhone 6s 的两倍。如果您提及您正在测试的设备以及您的目标性能特征是什么，人们可能会提供更多帮助。

@warrenm 我已经添加了细节。输入：128 * 51200 * 2 * sizeof(float)，输出：130806 * sizeof(float) 每帧。

我想知道使用 mmap + mlock 是否会有所帮助。

@nielsbot 我应该指出我也尝试过 let metalBuffer = self.metalDevice.newBufferWithLength(byteCount, options: .StorageModeShared)。这对性能没有影响。人们会认为至少 Apple API 调用会知道如何才能获得最佳性能。

很公平。另外，不要太迂腐，但 ios 上的页面并不总是 4k。

【参考方案1】：

减小类型的大小

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

Objective-C 优于 Swift

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

测试难度

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

【讨论】：

如果您不仅对iOS 设备感兴趣，您仍然可以在操场上运行测试，但目标是OS X。

@Marius 我只对 iOS 感兴趣，但还是不错的

【参考方案2】：

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

【讨论】：

这很有趣。感谢您的提示！

如果您喜欢 huffman GPU 解码器，我刚刚上传了基于 Rice 的 Metal GPU 解码器的源代码，它的压缩效果明显更好，执行速度几乎是 huffman 版本的 2 倍。在github上：github.com/mdejong/MetalRice