《CUDA编程》5.获得GPU加速的关键

从本章起，将关注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));

1. 首先定义两个cudaEvent_t的变量start和stop，并用cudaEventCreate()初始化。
2. 在需要计时的代码块运行之前，把start传入cudaEventRecord()。
3. 若是在TCC驱动模式的GPU中，cudaEventQuery(start);可以省略；若是处于WDDM驱动模式中，则必须保留。关于这两种模式，后面会讨论。
4. 在需要计时的代码块运行之后，把stop传入cudaEventRecord()。
5. 使用cudaEventSynchronize(stop)让主机等待事件stop被记录完毕
6. 调用 cudaEventElapsedTime() 函数计算 start 和 stop 这两个事件之间的时间差（单位是 ms）并输出到屏幕。
7. 调用 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上进行运算的数据规模。

1. 左图N在3次方和4次方时，耗时差距不大，是因为这两个时候，GPU的线程还有空闲，即所有数据都有线程在处理。从5次方开始，因为数据规模以及超过线程数量了，所以需要排队等待计算，故而随着N的增加，耗时成比例增加
2. 有图是和CPU相比，GPU的运算加速比。可以发现在3次方和4次方时，加速比增大，因为GPU有大量的线程可以并行，远远比CPU运算快。在达到5次方时，线程以及利用完毕，也得排队等待计算，所以加速比几乎不变

3 总结

在编写CUDA程序时，一定要做到以下三点：

1. 减少主机和设备之间的数据传输时间占比、也要减少数据传输次数
2. 提高核函数的算术强度
3. 增大核函数的并行规模

附：下面给出CUDA自带的数学函数库网站

http://docs.nvidia.com/cuda/cuda-math-api
包含幂函数、三角函数、指数函数、对数函数等，在编写代码时，要注意单精度和双精度的使用范围，例如有的计算精度不高的计算可以使用单精度，可以大大提升CUDA程序性能