cuda - GPU内存带宽理论与实际

Question

作为在 GPU 上运行的算法分析的一部分，我觉得我正在达到内存带宽。

我有几个复杂的内核执行一些复杂的操作(稀疏矩阵乘法、归约等)和一些非常简单的操作，当我计算读取/写入的总数据时，似乎所有(重要的)都达到了 ~79GB/s 带宽壁垒对于它们中的每一个，无论它们的复杂性如何，而理论 GPU 带宽为 112GB/s (nVidia GTX 960)

数据集非常大，在大约 10,000,000 个浮点条目的向量上运行，因此我从 clGetEventProfilingInfo 获得了很好的测量/统计数据，介于 COMMAND_START 和 COMMAND_END。在算法运行期间，所有数据都保留在 GPU 内存中，因此几乎没有主机/设备内存传输(也不是通过分析计数器测量的)

即使对于解决 x=x+alpha*b 的非常简单的内核(见下文)，其中 x 和 b 是约 10,000,000 个条目的巨大向量，我也没有接近理论带宽 (112GB/s) 而是在最大带宽 (~79GB/s) 的 ~70% 上运行

__kernel void add_vectors(int N,__global float *x,__global float const *b,float factor)
{
    int gid = get_global_id(0);
    if(gid < N)
        x[gid]+=b[gid]*factor;
}

我计算这个特定内核每次运行的数据传输为 N * (2 + 1) * 4:

• N - 矢量大小 = ~10,000,000
• 每个向量条目 2 次加载和 1 次存储
• 4 表示 float 大小

我预计对于这样一个简单的内核，我需要更接近带宽限制，我错过了什么？

P.S.:我从相同算法的 CUDA 实现中得到了相似的数字

Answer 1

我认为评估您是否已达到峰值带宽的更现实的方法是将您获得的内容与简单的 D2D 副本进行比较。

例如，你的内核读取 x 和 b 一次，写入 x 一次，因此执行时间的上限应该是从 b 复制到 x 一次的 1.5 倍。如果你发现时间比 1.5 倍高很多，这意味着你可能还有改进的空间。在这个内核中，工作非常简单，以至于开销(启动和结束函数、计算索引等)可能会限制性能。如果这是一个问题，您可能会发现使用网格跨步循环增加每个线程的工作量会有所帮助。

https://devblogs.nvidia.com/parallelforall/cuda-pro-tip-write-flexible-kernels-grid-stride-loops/

至于理论带宽，至少你应该考虑ECC如果启用的话开销。