从本章起，将关注CDUA程序的性能，即执行速度
1 用CUDA事件计时
在前几章中，使用的是C++的<time.h>库进行程序运行计时，CUDA也提供了一种基于CUDA event的计时方式，用来给一段CUDA代码进行计时，这里只介绍基于cudaEvent_t的计时方式，下面是一代码框架：
#include <cuda_runtime.h>
#include "error_check.cuh"
cudaEvent_t start, stop;
CHECK(cudaEventCreate(&start));
CHECK(cudaEventCreate(&stop));
CHECK(cudaEventRecord(start));
cudaEventQuery(start);//不能使用CHECK，因为可能返回cudaErrorNotReady，但并不是代码报错
/*需要计时的代码块*/
CHECK(cudaEventRecord(stop));
CHECK(cudaEventSynchronize(stop));
float time;
CHECK(cudaEventElapsedTime(&time, start, stop));
printf("time: %f ms\n", time);
CHECK(cudaEventDestroy(start));
CHECK(cudaEventDestroy(stop));
- 首先定义两个
cudaEvent_t的变量start和stop，并用cudaEventCreate()初始化。 - 在需要计时的代码块运行之前，把
start传入cudaEventRecord()。 - 若是在TCC驱动模式的GPU中，
cudaEventQuery(start);可以省略；若是处于WDDM驱动模式中，则必须保留。关于这两种模式，后面会讨论。 - 在需要计时的代码块运行之后，把
stop传入cudaEventRecord()。 - 使用
cudaEventSynchronize(stop)让主机等待事件stop被记录完毕 - 调用
cudaEventElapsedTime()函数计算start和stop这两个事件之间的时间差（单位是 ms）并输出到屏幕。 - 调用
cudaEventDestroy()函数销毁start和stop这两个 CUDA 事件。
1.1 举例说明
下面是一段利用cudaEvent_t进行计时的代码，在调用核函数add前后进行计时，注意，为了能够进行单精度和双精度的比较，需要定义一个宏变量，以便在编译时选择精度
#include <cuda.h>
#include <cuda_runtime.h>
#include <math.h>
#include <stdio.h>
#include "error_check.cuh"
#ifdef USE_DP
typedef double real;
const real EPSILON = 1.0e-15;
#else
typedef float real;
const real EPSILON = 1.0e-6f;
#endif
const real EPS = 1.0e-15;
const real a = 1.23;
const real b = 2.34;
const real c = 3.57;
// 希望 add 函数在 GPU 上执行
__global__ void add(const real* x, const real* y, real* z);
void check(const real* z, const int N);
int main(void) {
const int N = 100000000; // 定义数组的长度为 10 的 8 次方
const int M = sizeof(real) * N; // 每个数组所需的字节数
// 分配host内存
real* h_x = (real*)malloc(M);
real* h_y = (real*)malloc(M);
real* h_z = (real*)malloc(M);
for (int n = 0; n < N; ++n) {
h_x[n] = a;
h_y[n] = b;
}
//分配device内存
real* d_x, * d_y, * d_z;
CHECK(cudaMalloc((void**)&d_x, M));
CHECK(cudaMalloc((void**)&d_y, M));
CHECK(cudaMalloc((void**)&d_z, M));
// 将数据从主机复制到设备上
CHECK(cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice));
CHECK(cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice));
const int block_size = 128;
// 计算网格尺寸，确保所有元素都能被处理
const int grid_size = (N + block_size - 1) / block_size;
cudaEvent_t start, stop;
CHECK(cudaEventCreate(&start));
CHECK(cudaEventCreate(&stop));
CHECK(cudaEventRecord(start));
// 调用内核函数在设备中进行计算
add << <grid_size, block_size >> > (d_x, d_y, d_z);
CHECK(cudaEventRecord(stop));
CHECK(cudaEventSynchronize(stop));
// 将计算结果从设备复制回主机
CHECK(cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost));
check(h_z, N);
float elapsed_time;
CHECK(cudaEventElapsedTime(&elapsed_time, start, stop));
printf("Elapsed time: %f ms\n", elapsed_time);
// 释放内存
CHECK(cudaEventDestroy(start));
CHECK(cudaEventDestroy(stop));
free(h_x);
free(h_y);
free(h_z);
CHECK(cudaFree(d_x));
CHECK(cudaFree(d_y));
CHECK(cudaFree(d_z));
return 0;
}
__global__ void add(const real* x, const real* y, real* z) {
const int n = blockIdx.x * blockDim.x + threadIdx.x;
z[n] = x[n] + y[n];
}
void check(const real* z, const int N) {
bool has_error = false;
for (int n = 0; n < N; ++n) {
if (fabs(z[n] - c) > EPS) {
has_error = true;
}
}
printf("Has error: %d\n", has_error);
}
①单精度编译过程和输出结果
nvcc -o singel_cuda addFunction.cu -arch=sm_75
运行后输出如下：
②双精度编译过程和输出结果
nvcc -DUSE_DP -o double_cuda addFunction.cu -arch=sm_75
运行后输出如下：
观察结果，我们发现单精度的运行时间是10.624ms；双精度的运行时间是21.490ms
1.2 该计算任务并不适合使用GPU进行加速
我们把1.1代码中的数据复制步骤也加入到计时当中，观察耗时情况：
①单精度输出结果
①双精度输出结果
观察发现，核函数的运行时间连整体运行时间的10%都没有，若是算上CPU和GPU之间的传输时间，把该程序放入GPU中运算的性能，可能还不如直接在CPU上运行。
这里可以使用CUDA自带的nvprof工具对程序进行性能分析：
nvprof .\singel_cuda.exe
输出结果如下：
根据分析结果可以得出
- Host-to-Device (HtoD) 内存复制：
占总 GPU 时间的 58.98%，耗时 127.74 毫秒。 - Device-to-Host (DtoH) 内存复制：
占总 GPU 时间的 36.13%，耗时 78.259 毫秒。 - CUDA 内核函数 add 的执行：
占总 GPU 时间的 4.89%，耗时 10.590 毫秒
这意味着在总的 GPU 活动时间中，大约 95.11% 的时间都花在了内存复制上，而只有 4.89% 的时间用于实际的计算。所以这样的任务，其实是不适合使用GPU进行 “加速” 的，那么什么任务才能真正的发挥GPU加速能力呢？
2 影响GPU加速的关键因素
2.1 数据传输比例
从上一个例子我们可以得出，如果一任务仅仅是计算两个数组的和，那么用GPU可能比用CPU还慢，因为花在CPU与GPU之间传输的时间比计算时间还要多太多。
所以一个适合使用GPU加速的任务，一定是数据传输占比时间少的任务，尽量让一些操作在GPU中完成，避免过多数据经过PCIe传输，例如做10000次数组相加，只在开头和结尾进行数据传输（H to D/D to H)
下面是把上面代码的add操作重复1000次：
// 调用内核函数在设备中进行计算 1000 次
for (int i = 0; i < 1000; ++i) {
add << <grid_size, block_size >> > (d_x, d_y, d_z);
}
运行性能分析，结果输出如下：
性能分析结果显示，add函数运行耗时占比97.96%，数据传输占比是2.04%，所以这样的任务就适合使用GPU加速
2.2 算术强度（arithmetic intensity）
算术强度： 计算过程中浮点运算次数与读写内存字节数的比例。
在上述例子中，我们只进行了加法运算，接下来我们修改核函数，进行一些更复杂的数学运算，代码如下：
__global__ void add(const real* x, const real* y, real* z) {
const int n = blockIdx.x * blockDim.x + threadIdx.x;
real x_val = x[n];
real y_val = y[n];
// 复杂的数学运算
real result = sin(x_val) + cos(y_val) + exp(x_val * y_val) / (1.0 + x_val * y_val);
result += log10(x_val + 1.0) * sqrt(y_val);
// 最终结果
z[n] = result;
}
性能分析结果如下：
性能分析结果显示，运算时间占比为31.71%，传输数据时间占比为68.29%，比之前只做一次加法运算的操作更适合用GPU加速（之前运算时间占比是4.89%）
2.3 并行规模
并行规模： 指的是在GPU上同时执行的线程数量。
因为GPU上的线程是可以并行执行的，在设计核函数时，尽量让设备中的所有线程都要参与到计算之中，可以最大程度的加速CUDA程序，下面是两幅图：
N是指放到GPU上进行运算的数据规模。
- 左图N在3次方和4次方时，耗时差距不大，是因为这两个时候，GPU的线程还有空闲，即所有数据都有线程在处理。从5次方开始，因为数据规模以及超过线程数量了，所以需要排队等待计算，故而随着N的增加，耗时成比例增加
- 有图是和CPU相比，GPU的运算加速比。可以发现在3次方和4次方时，加速比增大，因为GPU有大量的线程可以并行，远远比CPU运算快。在达到5次方时，线程以及利用完毕，也得排队等待计算，所以加速比几乎不变
3 总结
在编写CUDA程序时，一定要做到以下三点：
- 减少主机和设备之间的数据传输时间占比、也要减少数据传输次数
- 提高核函数的算术强度
- 增大核函数的并行规模
附：下面给出CUDA自带的数学函数库网站
http://docs.nvidia.com/cuda/cuda-math-api
包含幂函数、三角函数、指数函数、对数函数等，在编写代码时，要注意单精度和双精度的使用范围，例如有的计算精度不高的计算可以使用单精度，可以大大提升CUDA程序性能
