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)，而是在~70% 的最大值 (~79GB/s)

__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 for sizeof 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的开销。