作者:翟天保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 forfor (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;}
测试效果
事件计时方法有精确度更高的计时效果,便于在调试和优化代码时使用。
如果函数有什么可以改进完善的地方,非常欢迎大家指出,一同进步何乐而不为呢~
如果文章帮助到你了,可以点个赞让我知道,我会很快乐~加油!