CUDA-事件计时方法cudaEventElapsedTime

已于 2024-09-21 16:28:53 修改
阅读量1.1k 收藏 10
点赞数 6
CC 4.0 BY-SA版权
分类专栏： CUDA 文章标签： 计时 cuda
于 2024-09-21 16:27:54 首次发布
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/zhaitianbao/article/details/142376957

作者：翟天保Steven
版权声明：著作权归作者所有，商业转载请联系作者获得授权，非商业转载请注明出处

实现原理

       事件计时方法是CUDA编程中一种用于测量GPU内核执行时间的高效手段。其基本步骤包括：

  1. 创建事件： 使用 cudaEvent_t 类型的变量来创建事件，通常有两个事件：一个用于记录开始时间，另一个用于记录结束时间。

  2. 记录开始事件： 在调用内核函数之前，使用 cudaEventRecord(start, 0) 记录开始时间。这个事件标记了内核执行的起始点。

  3. 调用GPU内核： 执行CUDA内核函数，例如使用kernel<<<gridSize, blockSize>>>(...)的方式。

  4. 记录结束事件： 内核执行完毕后，使用 cudaEventRecord(end, 0) 记录结束时间。这个事件标记了内核执行的结束点。

  5. 同步事件： 使用 cudaEventSynchronize(end) 确保CPU等待GPU完成内核执行，以确保在计算时间时所有操作都已结束。

  6. 计算耗时： 使用 cudaEventElapsedTime(&spendtime, start, end) 计算两个事件之间的时间差，将结果存储在指定变量中，通常以毫秒为单位。

  7. 输出结果： 将测量的时间输出，帮助分析和优化GPU内核的性能。

       该方法相比较一般的计时方法，如clock_t等，更适合在GPU场景计时，它有如下好处：

  1. 高精度： CUDA事件提供了高精度的计时，能够精确到毫秒级别，适合测量GPU内核的性能。

  2. 非阻塞： 使用事件记录的计时方式不会阻塞主机线程，允许CPU与GPU并行工作。这种非阻塞特性使得你可以在GPU执行的同时进行其他计算或处理。

  3. 自动同步： 通过调用 cudaEventSynchronize，你可以确保在计算时间时GPU的所有操作都已完成，避免了时间测量的不准确性。

  4. 易于使用： CUDA事件的使用相对简单，易于集成到现有的CUDA程序中，不需要复杂的编程或额外的库。

  5. 适用于多次测量： 可以很方便地进行多次测量，帮助你比较不同算法或参数设置的性能，进行性能调优。

  6. 支持并行计算： 你可以在多个内核调用中使用事件，监测不同内核的执行时间，进而优化内核的调度和资源利用。

  7. 便于调试和分析： 通过分析内核的执行时间，你可以识别瓶颈，帮助你更好地理解GPU的性能，并进行相应的优化。

       本文将通过一个之前分享过的中值滤波CUDA代码，进行事件计时方法的应用展示。

C++测试代码

Filter.h

#pragma once
#include <cuda_runtime.h>
#include <iostream>
#include <vector>
#include <opencv2/opencv.hpp>
#include <device_launch_parameters.h>

using namespace cv;
using namespace std;

#define TILE_WIDTH 16

// 预准备过程
void warmupCUDA();

// 中值滤波-CPU
cv::Mat filterMedian_CPU(cv::Mat input, int FilterWindowSize);

// 中值滤波-GPU
cv::Mat filterMedian_GPU(cv::Mat input, int FilterWindowSize);

Filter.cu

#include "Filter.h"

// 预准备过程
void warmupCUDA()
{
    float* dummy_data;
    cudaMalloc((void**)&dummy_data, sizeof(float));
    cudaFree(dummy_data);
}

// 中值滤波-CPU
cv::Mat filterMedian_CPU(cv::Mat input, int FilterWindowSize)
{
	int row = input.rows;
	int col = input.cols;

	// 预设输出
	cv::Mat output = input.clone();

	// 中值滤波
	int r = FilterWindowSize / 2;
#pragma omp parallel for
	for (int i = 0; i < row; ++i)
	{
		vector<uchar> datas;
		for (int j = 0; j < col; ++j)
		{
			// 卷积窗口边界限制，防止越界
			int ms = ((i - r) > 0) ? (i - r) : 0;
			int me = ((i + r) < (row - 1)) ? (i + r) : (row - 1);
			int ns = ((j - r) > 0) ? (j - r) : 0;
			int ne = ((j + r) < (col - 1)) ? (j + r) : (col - 1);
			// 求窗口内有效数据的中值
			datas.clear();
			for (int m = ms; m <= me; ++m)
			{
				for (int n = ns; n <= ne; ++n)
				{
					datas.push_back(input.at<uchar>(m, n));
				}
			}
			sort(datas.begin(), datas.end());
			output.at<uchar>(i, j) = datas[datas.size() / 2];
		}
	}

	return output;
}
// 中值滤波核函数
__global__ void medianFilter_CUDA(uchar* inputImage, uchar* outputImage, int width, int height, int windowSize)
{
	int row = blockIdx.y * blockDim.y + threadIdx.y;
	int col = blockIdx.x * blockDim.x + threadIdx.x;

	if (row < height && col < width)
	{
		// 参数预设
		uchar datas[100];
		int r = windowSize / 2;
		int ms = max(row - r, 0);
		int me = min(row + r, height - 1);
		int ns = max(col - r, 0);
		int ne = min(col + r, width - 1);
		// 赋值
		int count = 0;
		for (int m = ms; m <= me; ++m)
		{
			for (int n = ns; n <= ne; ++n)
			{
				datas[count++] = inputImage[m * width + n];
			}
		}
		// 选择排序
		for (int i = 0; i < count - 1; i++)
		{
			int minIndex = i;
			for (int j = i + 1; j < count; j++)
			{
				if (datas[j] < datas[minIndex])
				{
					minIndex = j;
				}
			}
			uchar temp = datas[i];
			datas[i] = datas[minIndex];
			datas[minIndex] = temp;
		}
		outputImage[row * width + col] = datas[count / 2];
	}
}
// 中值滤波-GPU
cv::Mat filterMedian_GPU(cv::Mat input, int FilterWindowSize)
{
	int row = input.rows;
	int col = input.cols;

	// 定义计时器
	float spendtime = 0.0f;
	cudaEvent_t start, end;
	cudaEventCreate(&start);
	cudaEventCreate(&end);

	// 分配GPU内存	
	uchar* d_inputImage, *d_outputImage;
	cudaMalloc(&d_inputImage, row * col * sizeof(uchar));
	cudaMalloc(&d_outputImage, row * col * sizeof(uchar));
	
	// 将输入图像数据从主机内存复制到GPU内存
	cudaMemcpy(d_inputImage, input.data, row * col * sizeof(uchar), cudaMemcpyHostToDevice);

	// 计算块和线程的大小
	dim3 blockSize(TILE_WIDTH, TILE_WIDTH);
	dim3 gridSize((col + blockSize.x - 1) / blockSize.x, (row + blockSize.y - 1) / blockSize.y);

	// 调用CUDA内核
	cudaEventRecord(start, 0);
	medianFilter_CUDA << <gridSize, blockSize >> > (d_inputImage, d_outputImage, col, row, FilterWindowSize);
	cudaEventRecord(end, 0);
	cudaEventSynchronize(end);
	cudaEventElapsedTime(&spendtime, start, end);
	cout << "medianFilter_CUDA spend time:" << spendtime << endl;

	// 将处理后的图像数据从GPU内存复制回主机内存
	cv::Mat output(row, col, CV_8UC1);
	cudaMemcpy(output.data, d_outputImage, row * col * sizeof(uchar), cudaMemcpyDeviceToHost);

	// 清理GPU内存
	cudaFree(d_inputImage);
	cudaFree(d_outputImage);

	return output;
}

main.cpp

#include "Filter.h"

void main()
{
    // 预准备
	warmupCUDA();

	cout << "medianFilter test begin." << endl;
	// 加载
	cv::Mat src = imread("test pic/test1.jpg", 0);
	int winSize = 9;
	cout << "filterWindowSize:" << winSize << endl;
	cout << "size: " << src.cols << " * " << src.rows << endl;

	// CPU版本
	clock_t s1, e1;
	s1 = clock();
	cv::Mat output1 = filterMedian_CPU(src, winSize);
	e1 = clock();
	cout << "CPU time:" << double(e1 - s1) / 1000 << endl;

	// GPU版本
	clock_t s2, e2;
	s2 = clock();
	cv::Mat output2 = filterMedian_GPU(src, winSize);
	e2 = clock();
	cout << "GPU time:" << double(e2 - s2) / 1000 << endl;

	// 检查
	int row = src.rows;
	int col = src.cols;
	bool flag = true;
	for (int i = 0; i < row; ++i)
	{
		for (int j = 0; j < col; ++j)
		{
			if (output1.at<uchar>(i, j) != output2.at<uchar>(i, j))
			{
				cout << "i:" << i << " j:" << j << endl;
				flag = false;
				break;
			}
		}
		if (!flag)
		{
			break;
		}
	}
	if (flag)
	{
		cout << "ok!" << endl;
	}
	else
	{
		cout << "error!" << endl;
	}

	// 查看输出
	cv::Mat test1 = output1.clone();
	cv::Mat test2 = output2.clone();

	cout << "medianFilter test end." << endl;
	
}

测试效果 

       事件计时方法有精确度更高的计时效果，便于在调试和优化代码时使用。

       如果函数有什么可以改进完善的地方，非常欢迎大家指出，一同进步何乐而不为呢~

       如果文章帮助到你了，可以点个赞让我知道，我会很快乐~加油！

CUDA —— 2.2、cuda - 计算核函数耗时 及 nvprof性能分析（附：完整代码）
07-11 291
CUDA —— 2.2、cuda - 计算核函数耗时（附：完整代码）
cuda示例asyncAPI--cuda事件GPU计时，和GPU、CPU执行中的使用
qq_43594926的博客
05-06 1394
此示例，说明了CUDA事件，在GPU和CPU同时执行时，计算GPU 的执行时间，即在GPU上程序运行的时间。 事件被插入到cuda调用流中，由于CUDA流调用，是异步的，即：在GPU执行时，CPU可以进行计算，CPU通过查询CUDA事件来判断GPU是否已经完成任务。 cuda核函数计时 cuda里面的程序是异步执行的， cpu将命令写入缓存区， GPU读取命令启动核函数，执行任务，返回结果，cpu一般不会等待cuda函数结束，会直接去做其他的事情。 一般GPU会给CPU汇报执行的进度，命令缓冲区和同步信息
cuda中的计时
qq_45001063的博客
04-24 283
代码解析：第1行cudaEvent_t start, stop：定义两个CUDA事件类型（cudaEvent_t）的变量；第2、3行cudaEventCreate函数初始化定义的cudaEvent_t变量；第4行通过cudaEventRecord函数，在需要记时的代码块之前记录代表时间开始的事件；第5行cudaEventQuery函数在TCC驱动模式的GPU下可省略，但在处于WDDM驱动模式的GPU必须保留，因此，我们就一直保留这句函数即可。
【CUDA学习】计时方法
weixin_34405332的博客
06-04 368
问题描述:一般利用CUDA进行加速处理时,都需要测试CUDA程序的运行时间,来对比得到的加速效果. 解决方法: 1).GPU端计时,即设备端计时. 2).CPU端计时,即主机端计时. 设备端计时有两种不同的方不地,分别是调用clock()函数和使用CUDA API的事件管理功能. clock函数计时: 1).在内核函数中要测量的一段代码的开始和结束的位置分别调用一次clock函数,并将...
[转载]CUDA 统计时间
weixin_33681778的博客
11-16 234
原文地址：CUDA统计时间作者：handsomefriend 转载：http://blog.csdn.net/jdhanhua/article/details/4843653 <1>使用cutil.h中的函数unsigned int timer=0;//创建计时器cutCreateTimer(&timer);//开始计时cutStartTimer(ti...
CUDA使用Event进行程序计时
weixin_33834137的博客
03-01 384
GPGPU是众核设备，包含大量的计算单元，实现超高速的并行。 使用CUDA在nvidia显卡上面编程时，可以使用CUDA提供的Event进行程序计时。 当然，每种编程语言基本都提供了获取系统时间的函数，如C/C++/Java 程序计时功能函数 Event可以统计GPU上面某一个任务或者代码段的精确运行时间。 如下面的程序实例(CalTime.cu)： 1 #include&lt;s...
cuda计时
han_hhh的博客
10-07 343
cudaEvent_t start, stop; cudaEventCreate(&start); //创建开始事件 cudaEventCreate(&stop); //创建结束事件 cudaEventRecord( start, 0 ); //记录起始时间 // do some work on the GPU cudaEventRecord( stop, 0 ); //记录结束时间 cudaEventSynchronize( stop ); //
CUDA编程之CUDA Sample-1_Utilities-bandwidthTest
翊桐公羽的博客
07-10 3493
CUDA Samples 中 1_Utilities 文件夹里包含了一些实用工具和小型示例程序,它们通常用于支持和演示其他 CUDA 示例程序的功能。bandwidthTest这个sample是一个简单的测试程序,用于测量 GPU 的内存拷贝带宽以及 PCI-E 总线上的内存拷贝带宽。这个程序可以帮助开发人员了解 CUDA 应用程序在不同内存访问模式下的性能表现,从而优化应用程序的内存使用策略,提高整体性能。
CUDA编程之CUDA Sample-0_Introduction-asyncAPI
翊桐公羽的博客
06-06 1723
包含库：代码包含了必要的系统和CUDA相关的头文件。内核函数：代码定义了一个CUDA内核函数，该内核函数接受一个整数数组g_data和一个增量值inc_value，并将数组中的每个元素增加给定的值。正确性检查：代码定义了一个函数，通过将data数组中的值与预期值进行比较来检查输出的正确性。主函数main()选择最佳可用的CUDA设备。为一个大小为1600万个元素的整数数组分配主机和设备内存。创建CUDA事件句柄以测量GPU的执行时间。设置内核启动配置，每个块512个线程，并计算适当的块数。
cuda核函数运行时间计时--简单
qq_43594926的博客
03-30 1991
** CPU与GPU同步并行与异步并行 **并行与并发 并发：针对运行在单处理器上程序的性能 单CPU上，同一时刻只能有一个线程运行，即：把时间分成许多段，轮巡调用不同的线程，不断切换 ，让不同的线程运行；即为并发。微观上，不是同步执行的 就好像去吃流水席，一个人在某个时间只能吃到，上面漂下来的一小份菜，但是几道菜会穿插着漂下来，比如顺序是：糖醋排骨，辣子鸡，四喜丸子，炖猪蹄， 每次按着这个顺序往下流，每次的分量是一整盘菜的一点儿，那么整个流水席菜的顺序是：0.1份糖醋排骨，0.2份辣子鸡，0.1份四喜丸子
CUDA-事件计时
DLLPJY的专栏
11-05 1137
cudaEvent_t start, stop; cudaEventCreate(&start); cudaEventCreate(&stop);cudaEventDestroy(start); cudaEventDestroy(stop);cudaEventRecord(start, 0); for (int i = 0; i < 2; ++i) { cudaMemcpyAsync(in
CUDA学习——统计时间
渐行渐远
11-20 9325
在CUDA中统计运算时间，大致有三种方法： 使用cutil.h中的函数unsigned int timer=0;//创建计时器cutCreateTimer(&timer);//开始计时cutStartTimer(timer);{  //统计的代码段  …………}//停止计时cutStopTimer(timer);//获得从开始计时到停止之间的时间cutGetTimerValue( time
CUDA编程【3】计时
qq_42681630的博客
01-03 1033
创建事件。：销毁事件。：在流中记录事件。：同步事件。：查询事件状态。：计算时间间隔。
GPU（七）CUDA事件计时与nvprof
NUCEMLS的博客
03-31 1116
CUDA计时与nvprof
CUDA 计时功能，记录GPU程序/函数耗时，cudaEventCreate，cudaEventRecord，cudaEventElapsedTime
爱串门的小马驹博客
12-09 808
为了测试GPU函数的耗时，可以使用 CUDA 提供的计时功能：cudaEventCreate, 和 cudaEventElapsedTime。这些函数可以帮助你测量某个 CUDA 操作（如设置设备）所花费的时间。
CUDA ---- Stream and Event
weixin_30902251的博客
09-20 766
Stream 一般来说，cuda c并行性表现在下面两个层面上： Kernel level Grid level 到目前为止，我们讨论的一直是kernel level的，也就是一个kernel或者一个task由许多thread并行的执行在GPU上。Stream的概念是相对于后者来说的，Grid level是指多个kernel在一个device上同时执行。 Stream和event...
CUDA计时
wanchaochaochao的博客
06-04 1153
CPU和GPU异步并发   CUDA内核程序的所有启动都是异步的，CPU通过将命令写入命令缓冲区，来请求启动内核，然后直接返回结果，而不检查GPU的执行进度。内存复制也可以选择异步方式，这使CPU和GPU可以并发执行，也使内存复制和内核处理可以并发执行。 CPU和GPU同步   在CUDA程序中，CPU端将命令写到命令缓冲区中，GPU端会依次读取命令缓冲区并执行任务，一般的程序中，GPU需要给CPU汇报任务的工作进度。命令缓冲区和“同步信息位置”都位于锁页主机内存上，所以CPU和GPU都能同时读写这些数据。
CUDA学习笔记 —— （六）程序计时器Events
Charles Ren's Tech Blog
04-05 1721
文章目录程序计时CPU计时器使用CUDA Events来计时Event来记录特定stream https://developer.nvidia.com/blog/how-implement-performance-metrics-cuda-cc/ 程序计时 CPU计时器 cudaMemcpy(d_x, x, N*sizeof(float), cudaMemcpyHostToDevice); cudaMemcpy(d_y, y, N*sizeof(float), cudaMemcpyHostToDevice)
CUDA的计时方法
Warren_Hoo的专栏
02-25 2266
CUDA Runtime API提供了用于计时的接口，可以用如下代码对在GPU上运行的CUDA程序进行计时： cudaEvent_t start, stop; cudaEventCreate(&start); cudaEventCreate(&stop); cudaEventRecord(start, 0); //需要计时的在GPU上运行的程序 cudaEventRecord(s
cuda事件
最新发布
03-19
### CUDA Events 的基本信息 CUDA 提供了一组用于时间测量和同步的事件接口，这些功能通过 `cudaEvent_t` 类型实现。以下是几个常用的 CUDA Event API 及其作用： - **`cudaEventCreate(cudaEvent_t *event)`**: 创建一个新的 CUDA 事件对象[^1]。 - **`cudaEventRecord(cudaEvent_t event, cudaStream_t stream=0)`**: 将指定事件记录到当前流中，在此之后可以用来标记某个操作完成的时间点。 - **`cudaEventDestroy(cudaEvent_t event)`**: 销毁之前创建的一个事件对象以释放资源。 - **`cudaEventElapsedTime(float *ms, cudaEvent_t start, cudaEvent_t end)`**: 计算两个事件之间经过的时间（单位为毫秒），这非常适用于性能分析场景下的精确计时需求。 下面是一个简单的例子展示如何利用上述API来计算数据传输所需时间: ```cpp #include <iostream> #include <cuda_runtime.h> int main() { float milliseconds = 0; // 定义并初始化设备指针与主机数组 (假设大小为N字节) size_t N = 1 << 20; // 1MB 数据量作为示例 char* h_data = new char[N]; char* d_data; // 分配显存空间 cudaMalloc((void**)&d_data, N); // 创建事件句柄 cudaEvent_t start_event, stop_event; cudaEventCreate(&start_event); cudaEventCreate(&stop_event); // 开始计时前先启动异步拷贝过程 cudaMemcpyAsync(d_data, h_data, N, cudaMemcpyHostToDevice, 0); // 插入开始结束标志位以便稍后统计耗时时长 cudaEventRecord(start_event, 0); cudaEventRecord(stop_event, 0); // 阻塞直到所有先前发生的命令都已完成执行(包括我们的copy指令),从而确保能够获取有效结果. cudaEventSynchronize(stop_event); // 获取实际花费时间 cudaEventElapsedTime(&milliseconds, start_event, stop_event); std::cout << "Time to transfer data: " << milliseconds << " ms." << std::endl; delete[] h_data; cudaFree(d_data); cudaEventDestroy(start_event); cudaEventDestroy(stop_event); return 0; } ``` 以上程序片段展示了怎样借助CUDA events 来评估内存复制操作的速度表现情况. ### 关于nvJPEG库的信息补充 虽然您询问的是关于CUDA events的内容，但是也提到了另一个重要的CUDA组件——nvJPEG 库。这是一个专门针对图像解码优化设计的高性能软件包，特别适合处理大规模批量图片文件转换任务。它允许开发者构建自定义解决方案来进行高效压缩或者无损存储等工作流程管理[^2]. ####
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包

打赏作者

翟天保Steven

你的鼓励将是我创作的最大动力

¥1 ¥2 ¥4 ¥6 ¥10 ¥20
扫码支付：¥1
获取中
扫码支付

您的余额不足，请更换扫码支付或充值

打赏作者

实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值