Cuda 中的有效带宽 [关闭]答案

【问题标题】：Effective Bandwidth in Cuda [closed]Cuda 中的有效带宽 [关闭]
【发布时间】：2019-04-25 17:10:33
【问题描述】：

在计算 Cuda 中的有效带宽时，我是否计算共享内存中的读/写次数。下面给出一个示例代码。

__global__ void kernel(float *a, float * b, float * c, int num){
    int i = threadIdx.x + blockIdx.x*blockDim.x;
    __shared__ a_shared[NUM];
    __shared__ b_shared[NUM];
    if (i < NUM){
       a_shared[i] = a[i];
       b_shared[i] = b[i]   
       c[i] = a_shared[i] + b_shared[i];
    }
}

【问题讨论】：

docs.nvidia.com/cuda/cuda-c-best-practices-guide/… 。这在 nvidia 文档的第 8.2.2 节中有记录
仔细阅读该部分向我表明，全局加载/存储在视图中，其中全局流量的后备存储是 DRAM。因此，共享活动不在视图中，也不应包括在内。

标签： c++ cuda gpu gpu-shared-memory

【解决方案1】：

关于您在上面的 cmets 中指出的最佳实践指南的the section，我会说答案是否定的，不应该包括共享流量。

我们怎么知道的？

计算有效带宽的主要目的是将其与理论带宽进行比较：

要准确测量性能，计算理论带宽和有效带宽很有用。当后者远低于前者时，设计或实现细节很可能会降低带宽，增加带宽应该是后续优化工作的首要目标。

然而，理论带宽计算仅包括到 DRAM 的全局内存流量：

使用这些数据项，NVIDIA Tesla M2090 的峰值理论内存带宽为 177.6 GB/s：

该数字是 DRAM 带宽。它不包括共享内存带宽。

分析器测量的参考均与全局内存流量有关，而不是共享内存：

请求的全局负载吞吐量

请求的全球商店吞吐量

全局负载吞吐量

全球商店吞吐量

DRAM 读取吞吐量

DRAM 写入吞吐量

在 CUDA 正式文档中我所知道的任何地方都没有记录计算理论共享内存带宽的方法，因此它不能包含在理论带宽计算中。因此，出于比较目的，包括对共享内存带宽的测量没有意义。

【讨论】：