NVIDIA CUDA2023春训营（八）CUDA 事件

Reference

• 【CUDA 基础】6.1 流和事件概述

CUDA 事件

**CUDA 事件（CUDA Event）**是 CUDA 流中应用程序跟踪进度的一个方式，其本质是流执行过程中的一个标记，可以检查正在执行的流的操作是否到达该点

可以将事件当成一个操作插入到流的众多操作中，当执行到该操作时，所做的工作就是设置 host 端的一个 flag 来标记是否完成

事件通常执行以下两个基本事务：

• 同步流执行
• 监控设备的进展

举例来说，可使用 CUDA 事件来对算法计时，在算法开始前设置一个事件，在算法结束后设置一个事件，两个事件的时间差即为算法执行时间

事件的资源分配与回收

CUDA 中封装了 CUDA 事件类型 cudaEvent_t，使用如下代码可声明一个事件

cudaEvent_t event;

在声明一个事件后，该事件无法使用，需要使用 cudaEventCreate() 来为其分配资源，在使用完毕后，使用 cudaEventDestroy() 回收资源

__host__ cudaError_t cudaEventCreate(cudaEvent_t* event)
__host__ __device__ cudaError_t cudaEventDestroy(cudaEvent_t event)

需要注意的是，与 CUDA 流的资源释放相似，如果进行事件资源回收时，相关操作未完成，则在操作完成后会立刻释放资源

事件时间间隔

cudaEventRecord() 用于将事件添加到流中，当不指定非空流时，默认添加进空流，函数原型如下：

__host__ __device__ cudaError_t cudaEventRecord(cudaEvent_t event, cudaStream_t stream = 0) 

cudaEventElapsedTime() 用于记录两个事件间的时间间隔，单位为毫秒，函数原型如下：

__host__ cudaError_t cudaEventElapsedTime(float* ms, cudaEvent_t start, cudaEvent_t end) 

两个事件 start 和 stop 不必关联到同一个流上，但需要注意的是，由于 cudaEventRecord() 是异步发生的，如果二者任一关联到非空流上，是无法保证度量出的时间间隔恰好是这两个事件的时间间隔，得到的时间间隔可能比期望的要大

如果只是想要查看 GPU 工作的时间间隔，将 start 和 stop 事件都默认添加进空流即可

事件的同步函数

在流中的事件主要是等待前面的操作完成后触发，与流同步函数类似，事件也具有阻塞和非阻塞两种同步函数

__host__ cudaError_t cudaEventSynchronize(cudaEvent_t event);
__host__ cudaError_t cudaEventQuery(cudaEvent_t event);

cudaEventSynchronize() 会强制 host 端阻塞等待，直到事件前的所有操作执行完成；cudaStreamQuery() 会检查事件前的操作是否全部完成，不会阻塞 host 端，如果事件前的所有操作都执行完成，那么会返回 cudaSuccess，否则返回 cudaErrorNotReady

实例

下述代码给出了记录事件时间间隔的实例

#include <stdio.h>
#define N 5
#define PI acos(-1.0)

__global__ void kernel(double *x, int n) {
    int index = threadIdx.x + blockIdx.x * blockDim.x;
    int stride = blockDim.x * gridDim.x;
    for (int i = index; i < n; i += stride) {
        x[i] = pow(PI,i);
    }
}

int main() {

    // 声明事件
    cudaEvent_t start, stop;

    // 分配事件资源
    cudaEventCreate(&start);
    cudaEventCreate(&stop);

    // 申请host锁定内存
    double *h_data;
    cudaMallocHost(&h_data, sizeof(double) * N);

    // 申请device内存
    double *d_data;
    cudaMalloc(&d_data, sizeof(double) * N);

    // 将start事件插入空流
    cudaEventRecord(start);

    // 执行核函数
    kernel<<<1, 64>>>(d_data, N);

    // 将stop事件插入空流
    cudaEventRecord(stop);

    // 阻塞host端直到事件完成
    cudaEventSynchronize(stop);

    // device数据传输到host
    cudaMemcpy(h_data, d_data, sizeof(double) * N, cudaMemcpyDeviceToHost);

    // 计算事件时间间隔
    float time;
    cudaEventElapsedTime(&time, start, stop);

    for (int i = 0; i < N; i++) {
        printf("%lf ", h_data[i]);
    }
    printf("\nGPU执行时间:%g ms\n", time);

    // 回收事件资源
    cudaEventDestroy(start);
    cudaEventDestroy(stop);
    // 释放host锁定内存
    cudaFreeHost(h_data);
    // 释放device显存
    cudaFree(d_data);
    
    return 0;
}