本文主要是介绍GPU（七）CUDA事件计时与nvprof，希望对大家解决编程问题提供一定的参考价值，需要的开发者们随着小编来一起学习吧！
CUDA runtime API文档参考：https://docs.nvidia.com/cuda/pdf/CUDA_Runtime_API.pdf
1、CUDA事件计时
不管是CPU程序，还是CUDA GPU程序，性能永远是一个重要的关注点，而评估一段程序的性能，耗时是一个非常直观的指标。假设我们要测试某个函数的耗时，一般会在该函数开始时计时一下，函数结束的时候计时一下，最终计算两个时间差。在普通程序中完全可以这么做，但是在CUDA GPU程序的host和device具有异构计算特性，如果把计时代码放在CPU处执行，则在第二次计时前，要做好同步操作。本章节讨论一种CUDA事件计时（也就是在device上计时的方式）来计算核函数执行时间的方法。
1.1 事件计时CUDA函数
本章节要用到的CUDA函数定义和功能如下：
cudaEventCreate
函数定义：__host__cudaError_t cudaEventCreate (cudaEvent_t *event)
函数功能：创建一个事件对象
cudaEventRecord
函数定义：__host____device__cudaError_t cudaEventRecord(cudaEvent_t event, cudaStream_t stream)
函数功能：记录一个事件
cudaEventQuery
函数定义：__host__cudaError_t cudaEventQuery (cudaEvent_t event)
函数功能：查询一个事件状态
cudaEventSynchronize
函数定义：__host__cudaError_t cudaEventSynchronize(cudaEvent_t event)
函数功能：等待一个事件完成
cudaEventElapsedTime
函数定义： __host__cudaError_t cudaEventElapsedTime (float *ms, cudaEvent_t start, cudaEvent_t end)
函数功能：计算两个事件的运行时间差
cudaEventDestroy
函数定义：__host____device__cudaError_t cudaEventDestroy(cudaEvent_t event)
函数功能：销毁一个事件对象
1.2 事件计时示例函数
基于上面的CUDA函数，我们写一个简单的CUDA程序：
// time.cu
#include<stdio.h>
#include<stdlib.h>
#include<time.h>const int g_buffer_size = 1024; // 数组大小
const int g_buffer_bytes = g_buffer_size*sizeof(float);
const int calc_count = 1000*1000; // 100w次计算
const long s2ns = 1000000000;
const long ms2ns = 1000000;void init_data(float *p_host){for (int i = 0; i < g_buffer_size; i++){p_host[i] = float(rand() % g_buffer_size);}
}__global__ void device_calc(float *p_device){const int idx = threadIdx.x + blockIdx.x*blockDim.x;p_device[idx] *= 1;
}void host_calc(float * p_host){for (int i = 0; i < g_buffer_size; i++){p_host[i] *= 1;}
}void host_time(float *p_host){struct timespec start, end;clock_gettime(CLOCK_REALTIME, &start);for (int i = 0; i < calc_count; i++){host_calc(p_host);}clock_gettime(CLOCK_REALTIME, &end);long start_us = start.tv_sec*s2ns+start.tv_nsec;long end_us = end.tv_sec*s2ns+end.tv_nsec;printf("host_time diff: %ld ms\n", (end_us-start_us)/ms2ns);
}void device_time(float *p_device){cudaEvent_t start, stop;cudaEventCreate(&start);cudaEventCreate(&stop);cudaEventRecord(start);cudaEventQuery(start);for (int i = 0; i < calc_count; i++){device_calc<<<1, g_buffer_size>>>(p_device);}cudaEventRecord(stop);cudaEventSynchronize(stop);float elapsed_time;cudaEventElapsedTime(&elapsed_time, start, stop);printf("device_time: %g ms\n", elapsed_time);cudaEventDestroy(start);cudaEventDestroy(stop);
}int main(){// host数组初始化float *p_host = (float*)malloc(g_buffer_bytes);if (p_host == NULL) {printf("malloc failed");exit(-1);}memset(p_host, 0, g_buffer_bytes);// 初始化host数据init_data(p_host);// device数组初始化float *p_device;cudaMalloc((float**)&p_device, g_buffer_bytes);if (p_device == NULL) {printf("cudaMalloc failed");free(p_host);exit(-1);}cudaMemset(p_device, 0, g_buffer_bytes);// 拷贝host数据到devicecudaMemcpy(p_device, p_host, g_buffer_bytes, cudaMemcpyHostToDevice);// host上计算时间host_time(p_host);// device上计算时间device_time(p_device);// 释放内存free(p_host);cudaFree(p_device);return 0;
}
在上面的CUDA程序中，我们分别用CPU和GPU对一个长度为1024的float数组做了100W次简单计算，然后计算耗时，注意在GPU计算的计时处，我们用的是CUDA事件计时方法。编译运行：
$ nvcc time.cu -o time
$ ./time
host_time diff: 997 ms
device_time: 6058.68 ms
从运行结果来看，CPU比GPU快很多，这是因为CPU单核运算速度比GPU单核更快，以及数据通过PCIe从host到device比较耗时等原因，在处理少量数据时会更有优势。但是如果把数据量和计算复杂度提上来，GPU就有更大的发挥空间了。
2、nvprof
除了像上面那样显示的在程序中加入计时来分析程序的性能，NVIDIA还提供了一个叫做nvprof的工具，我们把上面的代码稍微修改一下来看看nvprof工具的用法：
// nvprof.cu
#include<stdio.h>
#include<stdlib.h>
#include<time.h>const int g_buffer_size = 1024; // 数组大小
const int g_buffer_bytes = g_buffer_size*sizeof(float);void init_data(float *p_host){for (int i = 0; i < g_buffer_size; i++){p_host[i] = float(rand() % g_buffer_size);}
}__global__ void device_calc(float *p_device){const int idx = threadIdx.x + blockIdx.x*blockDim.x;p_device[idx] *= 1;
}int main(){// host数组初始化float *p_host = (float*)malloc(g_buffer_bytes);if (p_host == NULL) {printf("malloc failed");exit(-1);}memset(p_host, 0, g_buffer_bytes);// 初始化host数据init_data(p_host);// device数组初始化float *p_device;cudaMalloc((float**)&p_device, g_buffer_bytes);if (p_device == NULL) {printf("cudaMalloc failed");free(p_host);exit(-1);}cudaMemset(p_device, 0, g_buffer_bytes);// 拷贝host数据到devicecudaMemcpy(p_device, p_host, g_buffer_bytes, cudaMemcpyHostToDevice);device_calc<<<1, g_buffer_size>>>(p_device);// 释放内存free(p_host);cudaFree(p_device);return 0;
}
编译CUDA程序生成可执行文件：
$ nvcc nvprof.cu -o prof
$ ll prof
-rwxr-xr-x 1 nuczzz nuczzz 989504 Mar 31 21:48 prof*
执行nsys nvprof ./prof得到如下结果，可以看出这段程序主要耗时是在cudaMalloc：
在一些低版本中，可以直接使用
nvprof ./prof，高版本需要加上nsys
$ nsys nvprof ./prof
WARNING: prof and any of its children processes will be profiled.Generating '/tmp/nsys-report-795c.qdstrm'
[1/7] [========================100%] report1.nsys-rep
[2/7] [========================100%] report1.sqlite
[3/7] Executing 'nvtx_sum' stats report
SKIPPED: /home/nuczzz/c/day4/report1.sqlite does not contain NV Tools Extension (NVTX) data.
[4/7] Executing 'cuda_api_sum' stats reportTime (%) Total Time (ns) Num Calls Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Name-------- --------------- --------- ----------- ----------- --------- --------- ----------- ----------------------99.8 180894728 1 180894728.0 180894728.0 180894728 180894728 0.0 cudaMalloc0.1 175601 1 175601.0 175601.0 175601 175601 0.0 cudaFree0.1 120538 1 120538.0 120538.0 120538 120538 0.0 cudaLaunchKernel0.0 62816 1 62816.0 62816.0 62816 62816 0.0 cudaMemcpy0.0 12792 1 12792.0 12792.0 12792 12792 0.0 cudaMemset0.0 804 1 804.0 804.0 804 804 0.0 cuModuleGetLoadingMode[5/7] Executing 'cuda_gpu_kern_sum' stats report
SKIPPED: /home/nuczzz/c/day4/report1.sqlite does not contain CUDA kernel data.
[6/7] Executing 'cuda_gpu_mem_time_sum' stats report
SKIPPED: /home/nuczzz/c/day4/report1.sqlite does not contain GPU memory data.
[7/7] Executing 'cuda_gpu_mem_size_sum' stats report
SKIPPED: /home/nuczzz/c/day4/report1.sqlite does not contain GPU memory data.
Generated:/home/nuczzz/c/day4/report1.nsys-rep/home/nuczzz/c/day4/report1.sqlite
微信公众号卡巴斯同步发布，欢迎大家关注。
这篇关于GPU（七）CUDA事件计时与nvprof的文章就介绍到这儿，希望我们推荐的文章对编程师们有所帮助！
