从本章起,将关注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程序性能