c++ - Cuda 中的有效带宽
问题描述
在计算 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];
}
}
解决方案
关于您在上述评论中指出的最佳实践指南部分,我认为答案是否定的,不应包括共享流量。
我们怎么知道呢?
- 计算有效带宽的主要目的是将其与理论带宽进行比较:
为了准确测量性能,计算理论带宽和有效带宽很有用。当后者远低于前者时,设计或实现细节很可能会降低带宽,增加带宽应该是后续优化工作的首要目标。
然而,理论带宽计算仅包括到 DRAM 的全局内存流量:
使用这些数据项,NVIDIA Tesla M2090 的峰值理论内存带宽为 177.6 GB/s:
该数字是 DRAM 带宽。它不包括共享内存带宽。
- 探查器测量的参考都与全局内存流量有关,而不是共享内存:
请求的全局负载吞吐量
请求的全球商店吞吐量
全局负载吞吐量
全球商店吞吐量
DRAM 读取吞吐量
DRAM 写入吞吐量
- 在 CUDA 正式文档中我所知道的任何地方都没有记录计算理论共享内存带宽的方法,因此它不能包含在理论带宽计算中。因此,出于比较目的,包括共享内存带宽的测量是没有意义的。
推荐阅读
- sql - 存储过程/查询在所有用例的 2% 中需要附加字段。最佳实践?
- sql-server - 首先在 EF Core 代码中将 DeleteBehavior 从 Cascade 更改为 Restrict
- java - “BufferedInputStream”与“配置大小的 FileInputStream”
- php - 按多列之和排序 laravel
- javascript - 我无法使用 POST 方法将值从 JS 传递到 PHP 表单
- java - JAVA Stream API 使用
- c++ - Android Studio 使用 Gcc 编译,如何将其更改为 G++?(使用 CMake)
- dialogflow-es - Dialogflow Webhook:如何解决 Webhook 状态代码 14 [不可用]?
- reactjs - 如何替换redux中的数组数据?
- python - Pandas 将多个相关列组合成单独的组