c++ - Cuda 中的有效带宽

Question

在计算 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];
    }
}

Answer 1

关于您在上述评论中指出的最佳实践指南部分，我认为答案是否定的，不应包括共享流量。

我们怎么知道呢？

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

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

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

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

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

2. 探查器测量的参考都与全局内存流量有关，而不是共享内存：

请求的全局负载吞吐量

请求的全球商店吞吐量

全局负载吞吐量

全球商店吞吐量

DRAM 读取吞吐量

DRAM 写入吞吐量

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