CUDA时间评估LZ之前都是使用CPU timer来评估CUDA代码的性能的但是这样做时候没有办法给出精确答案的除了其他因素外最常见的时间计算误差包括线程延迟开销和操作系统调度的开销。同时使用GPU测量的时间还取决于高精度Cpu计时器的可用性。在很多情况下主机在GPU内核运行时执行异步计算因此因此为了测量GPU内核消耗的时间CUDA提供了对应于event的API。CUDA event是在CUDA程序中指定位置记录GPU的时间戳在这个API中GPU记录了对应的时间戳消除了上面提到的两个问题。使用CUDA event有两个测量时间的步骤创建事件和记录事件最后计算中间的时间得到最后代码的整体性能。具体的代码块如下所示#includecuda_runtime_api.h#includestdio.hintmain(){// 一些数据准备......// 实例化CUDA eventcudaEvent_te_start,e_stop;//创建事件cudaEventCreate(e_start);cudaEventCreate(e_stop);//记录事件开始计算时间cudaEventRecord(e_start,0);...//一顿操作猛如虎...//记录结束时事件cudaEventRecord(e_stop,0);// 0 代表CUDA流0//等待事件同步后cudaEventSynchronize(e_stop);//计算对应的时间评估代码性能floatelapsedTime;cudaEventElapsedTime(elapsedTime,e_start,e_stop);}CUDA stream(CUDA 流)的一些理解使用多个block或者thread只是把一个任务进行并行化操作但是并没有体现任务的并行性。任务的并行性是指可以运行多个核函数最常见的如图像处理可以一个流做图像上下翻转一个流做随机crop还可以一个流做resize但是这个并不是那么灵活。CUDA流其实是在GPU执行中GPU操作队列具体顺序。这些功能包括内核功能内存复制操作和CUDA事件操作而这些操作添加队列的顺序其实就是对应操作执行的顺序。每个CUDA流可以看做一个单独的任务所以我们可以使用多个流来开启多个任务。#includecuda_runtime_api.h#includeiostream#includestdio.h#defineRANDOM(x)(rand()%x)#defineMAX10#defineM128// number of thread#defineN50000// Define number of elements in Array//Define kernel function for vector addition__global__voidgpuAdd(int*d_a,int*d_b,int*d_c){// getting block index of current kernelinttidthreadIdx.xblockIdx.x*blockDim.x;while(tidN){d_c[tid]d_a[tid]d_b[tid];tidblockDim.x*gridDim.x;}}intmain(){// define host arraysint*h_a,*h_b,*h_c;// define device pointers for stream0int*d_a0,*d_b0,*d_c0;// define device pointers for stream1int*d_a1,*d_b1,*d_c1;cudaStream_tstream0,stream1;cudaStreamCreate(stream0);cudaStreamCreate(stream1);cudaEvent_te_start,e_stop;cudaEventCreate(e_start);cudaEventCreate(e_stop);cudaEventRecord(e_start,0);// Allocate memory for host pointerscudaHostAlloc((void**)h_a,2*N*sizeof(int),cudaHostAllocDefault);cudaHostAlloc((void**)h_b,2*N*sizeof(int),cudaHostAllocDefault);cudaHostAlloc((void**)h_c,2*N*sizeof(int),cudaHostAllocDefault);// Allocate memory for device pointerscudaMalloc((void**)d_a0,N*sizeof(int));cudaMalloc((void**)d_b0,N*sizeof(int));cudaMalloc((void**)d_c0,N*sizeof(int));cudaMalloc((void**)d_a1,N*sizeof(int));cudaMalloc((void**)d_b1,N*sizeof(int));cudaMalloc((void**)d_c1,N*sizeof(int));for(inti0;iN*2;i){h_a[i]2*i*i;h_b[i]i;}cudaMemcpyAsync(d_a0,h_a,N*sizeof(int),cudaMemcpyHostToDevice,stream0);cudaMemcpyAsync(d_a1,h_aN,N*sizeof(int),cudaMemcpyHostToDevice,stream1);cudaMemcpyAsync(d_b0,h_b,N*sizeof(int),cudaMemcpyHostToDevice,stream0);cudaMemcpyAsync(d_b1,h_bN,N*sizeof(int),cudaMemcpyHostToDevice,stream1);// Kernel callgpuAdd512,512,0,stream0(d_a0,d_b0,d_c0);gpuAdd512,512,0,stream1(d_a1,d_b1,d_c1);// Copy result back to host memory from device memorycudaMemcpyAsync(h_c,d_c0,N*sizeof(int),cudaMemcpyDeviceToHost,stream0);cudaMemcpyAsync(h_cN,d_c1,N*sizeof(int),cudaMemcpyDeviceToHost,stream1);cudaDeviceSynchronize();cudaStreamSynchronize(stream0);cudaStreamSynchronize(stream1);cudaEventRecord(e_stop,0);cudaEventSynchronize(e_stop);floatelapsedTime;cudaEventElapsedTime(elapsedTime,e_start,e_stop);printf(Time to add %d numbers: %3.1f ms\n,2*N,elapsedTime);intCorrect1;printf(Vector addition on GPU: \n);// print result on consolefor(inti0;i2*N;i){if(h_a[i]h_b[i]!h_c[i]){Correct0;}}if(Correct1){printf(GPU has computed Sum Correctly\n);}else{printf(There is an Error in GPU computation\n);}// free up memorycudaFree(d_a0);cudaFree(d_b0);cudaFree(d_c0);cudaFree(d_a1);cudaFree(d_b1);cudaFree(d_c1);cudaFreeHost(h_a);cudaFreeHost(h_b);cudaFreeHost(h_c);return0;}附一张运行示意图最后运行结果如图所示参考书籍Hands-On GPU-Accelerated Computer Vision with OpenCV and CUDAPS每天起床第一件事就是看下疫情有没有缓解有些时候也只能干着急帮不上什么忙只能默默宅在家里不出门添乱了按照目前态势周六可能就破三万了心里慌慌的公司也是准备10号开始远程办公怕人员太多不便于管理的角度考虑这个做法还是值得赞赏的。但是觉得这个疫情情况界面可以再做好一些对每个区域的新增人数最好也有直观的表示而且是不是确诊病例绝大多数都是从疑似病例中确诊出来的那么新增的疑似病例如果分析一下可能昨天一天就增加了8000多例武汉加油湖北加油中国加油劲头上来了能行