GPU内存带宽理论值与实际值的差异

Question

GPU内存带宽理论值与实际值的差异

3

作为在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代表float类型的sizeof值

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

P.S.: 我在CUDA实现相同算法时得到类似的结果。

- Artyom

1个回答

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- kangshiyin · Accepted Answer

我认为评估是否达到峰值带宽的更现实方法是将其与简单的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，至少需要考虑其开销。