CUDA 程序的优化(2) 测量程序运行时间

您所在的位置:网站首页 怎么看程序运行记录 CUDA 程序的优化(2) 测量程序运行时间

CUDA 程序的优化(2) 测量程序运行时间

2024-07-11 19:39| 来源: 网络整理| 查看: 265

测量程序运行时间

本节将介绍如何准确地测量CUDA程序的运行时间。CUDA的内核程序运行时间可以在设备端测量,也可以在主机端测量。而CUDA API的运行时间则只能从主机端测量。无论是主机端测时还是设备端测时,最好都测量内核函数多次运行的时间,然后再除以运行次数以获得更加准确的结果。使用CUDA runtime API时,会在第一次调用runtime API函数时启动CUDA环境。为了避免将这一部分时间计入,最好在正式测时开始前先进行一次包含数据输入输出的计算,这样也可以使GPU从平时的节能模式进入工作状态,使测试结果更加可靠。

4.2.1设备端测时

设备端测时使用GPU中的计时器的时戳计时。实现设备端测时有两种不同的方法,分别是调用clock()函数和使用CUDA API的事件管理功能。

使用clock()函数计时,在内核函数中要测量的一段代码的开始和结束的位置分别调用一次clock()函数,并将结果记录下来。由于调用_synothreads()函数后,一个block中的所有thread需要的时间是相同的,因此只需要记录每个block执行需要的时间就行了,而不需要记录每个thread的时间。clock()函数的返回值的单位是GPU的时钟周期,需要除以GPU的运行频率才能得到以秒为单位的时间。这里测得的时间是一个block在GPU中上下文保持的时间,而不是实际执行需要的时间;每个block实际执行的时间一般要短于测得的结果。下面是一个使用clock函数测时的例子。

设备端代码:

#ifndef _CLOCK_KERNEL_H_ #define _CLOCK_KERNEL_H_ // 这段代码测量进行归约运算时每个 block 使用的时钟周期数,并将结果存储在显存中 __global__ static void timedReduction(const float * input, float * output, clock_t * timer) { // __shared__ float shared[2 * blockDim.x]; extern __shared__ float shared[]; const int tid = threadIdx.x; const int bid = blockIdx.x; //记录测时开始时的时戳 if (tid == 0) timer[bid] = clock(); // Copy input. shared[tid] = input[tid]; shared[tid + blockDim.x] = input[tid + blockDim.x]; // Perform reduction to find minimum. for(int d = blockDim.x; d > 0; d /= 2) { __syncthreads(); if (tid


【本文地址】


今日新闻


推荐新闻


CopyRight 2018-2019 办公设备维修网 版权所有 豫ICP备15022753号-3