CUDA 性能编程指南

https://developer.nvidia.com/zh-cn/blog/cuda-performance-guide-cn/
英伟达官方提供的文档：
2. CUDA C++ 最佳实践指南

1. 计时

cuda核函数执行时间可以使用CPU计时器或GPU计时器实现，任何 CPU 计时器都可用于测量 CUDA 调用或内核执行的运行时间。需要注意的时，很多cuda指令是异步的，也就是说，它们在完成工作之前将控制权交还给调用的 CPU 线程。所有内核启动都是异步的，名称上有后缀的内存复制函数也是如此。因此，为了准确测量特定调用或 CUDA 调用序列的运行时间，有必要通过在启动和停止 CPU 计时器之前立即调用来将 CPU 线程与 GPU 同步。阻止调用 CPU 线程，直到线程之前发出的所有 CUDA 调用都完成。[2]
CPU 到 GPU 同步点意味着 GPU 的处理管道中可能存在停滞，因此应谨慎使用，以最大程度地减少其性能影响.
CUDA 事件 API 提供用于创建和销毁事件、记录事件（包括时间戳）以及将时间戳差异转换为以毫秒为单位的浮点值的调用。使用cuda的事件进行计时，参考代码如下：

cudaEvent_t start, stop;
float time;

cudaEventCreate(&start);
cudaEventCreate(&stop);

cudaEventRecord( start, 0 );
kernel<<<grid,threads>>> ( d_odata, d_idata, size_x, size_y,
                          NUM_REPS);
cudaEventRecord( stop, 0 );
cudaEventSynchronize( stop );

cudaEventElapsedTime( &time, start, stop );
cudaEventDestroy( start );
cudaEventDestroy( stop );

时间统计结果毫秒为单位表示，分辨率约为半微秒。

2. 带宽

带宽 - 数据传输的速率 - 是影响性能的最重要门控因素之一。几乎所有对代码的更改都应该在它们如何影响带宽的上下文中进行。如本指南的内存优化中所述，带宽可能会受到存储数据的内存选择、数据布局方式和访问顺序以及其他因素的显著影响。[2]