4

NVIDIA CUDA2023春训营（八）CUDA 事件

 1 year ago
source link: https://alex-mcavoy.github.io/nvidia/cuda-spring-bootcamp/196d4211.html
Go to the source link to view the article. You can view the picture content, updated content and better typesetting reading experience. If the link is broken, please click the button below to view the snapshot at that time.
neoserver,ios ssh client

Reference

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;
}

report

Recommend

About Joyk


Aggregate valuable and interesting links.
Joyk means Joy of geeK