GPU内存带宽理论与实际

Question

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

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

数据集非常大，在约 10,000,000 个浮点条目的向量上运行，所以我在 COMMAND_START 和 COMMAND_END 之间从 clGetEventProfilingInfo 得到很好的 measurements/statistics。在算法运行期间，所有数据都保留在 GPU 内存中，因此几乎没有 host/device 内存传输（也不是通过分析计数器测量的）

即使对于解决 x=x+alpha*b 的非常简单的内核（见下文），其中 x 和 b 是约 10,000,000 个条目的巨大向量，我也没有接近理论带宽 (112GB/s) 但是而是运行ning 在最大值的 ~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 倍高很多，则表示您可能还有 space 需要改进。在这个内核中，工作非常简单，以至于开销（启动和结束函数、计算索引等）可能会限制性能。如果这是一个问题，您可能会发现使用网格跨步循环增加每个线程的工作量会有所帮助。

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

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

GPU内存带宽理论与实际

GPU Memory bandwidth theoretical vs practical

cuda

opencl

linear-algebra

gpgpu

bandwidth