CUDA-事件计时方法cudaEventElapsedTime

原创已于 2024-09-21 16:28:53 修改 · 1.1k 阅读

10 ·

CC 4.0 BY-SA版权

文章标签：

#计时 #cuda

于 2024-09-21 16:27:54 首次发布

CUDA 专栏收录该内容

11 篇文章

订阅专栏

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

实现原理

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

创建事件：使用 cudaEvent_t 类型的变量来创建事件，通常有两个事件：一个用于记录开始时间，另一个用于记录结束时间。
记录开始事件：在调用内核函数之前，使用 cudaEventRecord(start, 0) 记录开始时间。这个事件标记了内核执行的起始点。
调用GPU内核：执行CUDA内核函数，例如使用kernel<<<gridSize, blockSize>>>(...)的方式。
记录结束事件：内核执行完毕后，使用 cudaEventRecord(end, 0) 记录结束时间。这个事件标记了内核执行的结束点。
同步事件：使用 cudaEventSynchronize(end) 确保CPU等待GPU完成内核执行，以确保在计算时间时所有操作都已结束。
计算耗时：使用 cudaEventElapsedTime(&spendtime, start, end) 计算两个事件之间的时间差，将结果存储在指定变量中，通常以毫秒为单位。
输出结果：将测量的时间输出，帮助分析和优化GPU内核的性能。

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

高精度： CUDA事件提供了高精度的计时，能够精确到毫秒级别，适合测量GPU内核的性能。
非阻塞：使用事件记录的计时方式不会阻塞主机线程，允许CPU与GPU并行工作。这种非阻塞特性使得你可以在GPU执行的同时进行其他计算或处理。
自动同步：通过调用 cudaEventSynchronize，你可以确保在计算时间时GPU的所有操作都已完成，避免了时间测量的不准确性。
易于使用： CUDA事件的使用相对简单，易于集成到现有的CUDA程序中，不需要复杂的编程或额外的库。
适用于多次测量：可以很方便地进行多次测量，帮助你比较不同算法或参数设置的性能，进行性能调优。
支持并行计算：你可以在多个内核调用中使用事件，监测不同内核的执行时间，进而优化内核的调度和资源利用。
便于调试和分析：通过分析内核的执行时间，你可以识别瓶颈，帮助你更好地理解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;
	
}