MindSpore:CUDA编程（五）Event

本文主要是介绍MindSpore:CUDA编程（五）Event，对大家解决编程问题具有一定的参考价值，需要的程序猿们随着小编来一起学习吧！

Event是CUDA中的事件，用于分析、检测CUDA程序中的错误。

一般我们会定义一个宏：

#pragma once
#include <stdio.h>
 
#define CHECK(call)                                   \
do                                                    \
{                                                     \
    const cudaError_t error_code = call;              \
    if (error_code != cudaSuccess)                    \
    {                                                 \
        printf("CUDA Error:\n");                      \
        printf("    File:       %s\n", __FILE__);     \
        printf("    Line:       %d\n", __LINE__);     \
        printf("    Error code: %d\n", error_code);   \
        printf("    Error text: %s\n",                \
            cudaGetErrorString(error_code));          \
        exit(1);                                      \
    }                                                 \
} while (0)

并在适当的位置使用这个宏来打印CUDA的错误日志。#pragma once, 不要放在源代码文件里，这个一般只放在头文件里的。（防止头文件被引入多次）

Event的调用有以下内容：

具体的顺序如下：

（1）声明Event（这里以计算核函数运行时间前后的start Event和stop Event为例）

cudaEvent_t start, stop;

（2）创建Event

CHECK(cudaEventCreate(&start));
CHECK(cudaEventCreate(&stop));

（3）添加Event（在合适的地方）

cudaEventRecord(start);
cudaEventRecord(stop);

（4）等待Event完成

（a）非堵塞方式——可以用于一些不需要等待的处理

cudaEventQuery(start);

（b）堵塞方式——可以用于执行核函数后等待核函数执行完毕后的处理

cudaEventSynchronize(stop);

（5）计算两个Event间隔时间

CHECK(cudaEventElapsedTime(&elapsed_time, start, stop));

（6）销毁Event

CHECK(cudaEventDestroy(start));
CHECK(cudaEventDestroy(stop));

以上次介绍的矩阵乘为例，完整的代码如下：

#pragma once
#include <stdio.h>
 
#define CHECK(call)                                   \
do                                                    \
{                                                     \
    const cudaError_t error_code = call;              \
    if (error_code != cudaSuccess)                    \
    {                                                 \
        printf("CUDA Error:\n");                      \
        printf("    File:       %s\n", __FILE__);     \
        printf("    Line:       %d\n", __LINE__);     \
        printf("    Error code: %d\n", error_code);   \
        printf("    Error text: %s\n",                \
            cudaGetErrorString(error_code));          \
        exit(1);                                      \
    }                                                 \
} while (0)
 
#include <stdio.h>
#include <math.h>
#include "error.cuh"
 
#define BLOCK_SIZE 32
 
__global__ void gpu_matrix_mult(int *a,int *b, int *c, int m, int n, int k)
{ 
    int row = blockIdx.y * blockDim.y + threadIdx.y; 
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int sum = 0;
    if( col < k && row < m) 
    {
        for(int i = 0; i < n; i++) 
        {
            sum += a[row * n + i] * b[i * k + col];
        }
        c[row * k + col] = sum;
    }
} 
 
void cpu_matrix_mult(int *h_a, int *h_b, int *h_result, int m, int n, int k) {
    for (int i = 0; i < m; ++i) 
    {
        for (int j = 0; j < k; ++j) 
        {
            int tmp = 0.0;
            for (int h = 0; h < n; ++h) 
            {
                tmp += h_a[i * n + h] * h_b[h * k + j];
            }
            h_result[i * k + j] = tmp;
        }
    }
}
 
int main(int argc, char const *argv[])
{
    int m=100;
    int n=100;
    int k=100;
    
    //声明Event
    cudaEvent_t start, stop, stop2, stop3 , stop4 ;
    
    //创建Event
    CHECK(cudaEventCreate(&start));
    CHECK(cudaEventCreate(&stop));
    CHECK(cudaEventCreate(&stop2));
 
    int *h_a, *h_b, *h_c, *h_cc;
    CHECK(cudaMallocHost((void **) &h_a, sizeof(int)*m*n));
    CHECK(cudaMallocHost((void **) &h_b, sizeof(int)*n*k));
    CHECK(cudaMallocHost((void **) &h_c, sizeof(int)*m*k));
    CHECK(cudaMallocHost((void **) &h_cc, sizeof(int)*m*k));
 
    for (int i = 0; i < m; ++i) {
        for (int j = 0; j < n; ++j) {
            h_a[i * n + j] = rand() % 1024;
        }
    }
 
    for (int i = 0; i < n; ++i) {
        for (int j = 0; j < k; ++j) {
            h_b[i * k + j] = rand() % 1024;
        }
    }
 
    int *d_a, *d_b, *d_c;
    CHECK(cudaMalloc((void **) &d_a, sizeof(int)*m*n));
    CHECK(cudaMalloc((void **) &d_b, sizeof(int)*n*k));
    CHECK(cudaMalloc((void **) &d_c, sizeof(int)*m*k));
 
    // copy matrix A and B from host to device memory
    CHECK(cudaMemcpy(d_a, h_a, sizeof(int)*m*n, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(d_b, h_b, sizeof(int)*n*k, cudaMemcpyHostToDevice));
 
    unsigned int grid_rows = (m + BLOCK_SIZE - 1) / BLOCK_SIZE;
    unsigned int grid_cols = (k + BLOCK_SIZE - 1) / BLOCK_SIZE;
    dim3 dimGrid(grid_cols, grid_rows);
    dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
   
    //开始start Event
    cudaEventRecord(start);
    //非阻塞模式
    cudaEventQuery(start);
 
    //gpu_matrix_mult<<<dimGrid, dimBlock>>>(d_a, d_b, d_c, m, n, k);   
    
    gpu_matrix_mult_shared<<<dimGrid, dimBlock>>>(d_a, d_b, d_c, m, n, k);  
 
    //开始stop Event
    cudaEventRecord(stop);
    //由于要等待核函数执行完毕，所以选择阻塞模式
    cudaEventSynchronize(stop);
    
    //计算时间 stop-start
    float elapsed_time;
    CHECK(cudaEventElapsedTime(&elapsed_time, start, stop));
    printf("start-》stop：Time = %g ms.\n", elapsed_time);
 
    
    cudaMemcpy(h_c, d_c, (sizeof(int)*m*k), cudaMemcpyDeviceToHost);
    //cudaThreadSynchronize();
 
    //开始stop2 Event
    CHECK(cudaEventRecord(stop2));
    //非阻塞模式
    //CHECK(cudaEventSynchronize(stop2));
    cudaEventQuery(stop2);
 
    //计算时间 stop-stop2
    float elapsed_time2;
    cudaEventElapsedTime(&elapsed_time2, stop, stop2);
    printf("stop-》stop2：Time = %g ms.\n", elapsed_time2);
 
    //销毁Event
    CHECK(cudaEventDestroy(start));
    CHECK(cudaEventDestroy(stop));
    CHECK(cudaEventDestroy(stop2));
 
    //CPU函数计算
    cpu_matrix_mult(h_a, h_b, h_cc, m, n, k);
 
    int ok = 1;
    for (int i = 0; i < m; ++i)
    {
        for (int j = 0; j < k; ++j)
        {
            if(fabs(h_cc[i*k + j] - h_c[i*k + j])>(1.0e-10))
            {
                
                ok = 0;
            }
        }
    }
 
    if(ok)
    {
        printf("Pass!!!\n");
    }
    else
    {
        printf("Error!!!\n");
    }
 
    // free memory
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    cudaFreeHost(h_a);
    cudaFreeHost(h_b);
    cudaFreeHost(h_c);
    return 0;
}

在Quardo P1000的GPU上执行：

这里以矩阵乘为例，打印了调用矩阵乘核函数的时间，以及后面 cudaMemcpy的时间。

我们强行将 CHECK(cudaMemcpy(d_b, h_b, sizeof(int)*n*k, cudaMemcpyHostToDevice)); 改为 CHECK(cudaMemcpy(d_b, h_b, sizeof(int)*n*k*2, cudaMemcpyHostToDevice));  故意让其出界。

再重新编译，运行，看看效果：

系统会告诉你 这行有错：

这样就可以跟踪出CUDA调用中的错误。

这里需要总结一下张小白在调试CHECK过程中发现的几个问题：

（1）如果没有 CHECK(cudaEventCreate()) 就直接调用 cudaEventRecord() 或者执行后面的Event函数，会导致打印不了信息。张小白当时对于stop2这个event就犯了这个错，导致 stop->stop2的时间怎么都打不出来。

（2）对于 cudaEventQuery() 是不能加 CHECK的，如果加了反而会报错：

在上面的环境中，如果您这样写：

CHECK(cudaEventQuery(stop2));

编译执行就会出现以下错误：

cudaEventQuery的cudaErrorNotReady代表了事件还没发生（还没有被记录），不代表错误。

这篇关于MindSpore:CUDA编程（五）Event的文章就介绍到这儿，希望我们推荐的文章对大家有所帮助，也希望大家多多支持为之网！