事件計時

CUDA事件是直接在GPU上實現(xiàn)的，因此它們不適用于對同時包含設(shè)備代碼和主機代碼的混合代碼計時。

• cudaEventCreate 創(chuàng)建一個事件
• cudaEventRecord 記錄一個事件
• cudaEventElapsedTime 計算兩個事件之間經(jīng)歷的時間，第一個參數(shù)為某個浮點變量的地址，在這個參數(shù)中將包含兩次事件之間經(jīng)歷的時間，單位是毫秒

#include <stdio.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include<math.h>
#include <malloc.h> 
#include <opencv2/opencv.hpp>
#include <stdlib.h>#define BLOCK_SIZE 1//圖像卷積 GPU
__global__ void sobel_gpu(unsigned char* in, unsigned char* out, const int Height, const int Width)
{int x = blockDim.x * blockIdx.x + threadIdx.x;int y = blockDim.y + blockIdx.y + threadIdx.y;int index = y * Width + x;int Gx = 0;int Gy = 0;unsigned char x0, x1, x2, x3, x4, x5, x6, x7, x8;if (x>0 && x<(Width-1) && y>0 && y<(Height-1)){x0 = in[(y - 1)*Width + (x - 1)];x1 = in[(y - 1)*Width + (x)];x2 = in[(y - 1)*Width + (x + 1)];x3 = in[(y)*Width + (x - 1)];x5 = in[(y)*Width + (x + 1)];x6 = in[(y + 1)*Width + (x - 1)];x7 = in[(y + 1)*Width + (x)];x8 = in[(y + 1)*Width + (x + 1)];Gx = (x0 + 2 * x3 + x6) - (x2 + 2 * x5 + x8);Gy = (x0 + 2 * x1 + x2) - (x6 + 2 * x7 + x8);out[index] = (abs(Gx) + abs(Gy)) / 2;}
}int main()
{cv::Mat src;src = cv::imread("complete004.jpg");cv::Mat grayImg,gaussImg;cv::cvtColor(src, grayImg, cv::COLOR_BGR2GRAY);cv::GaussianBlur(grayImg, gaussImg, cv::Size(3,3), 0, 0, cv::BORDER_DEFAULT);int height = src.rows;int width = src.cols;//輸出圖像cv::Mat dst_gpu(height, width, CV_8UC1, cv::Scalar(0));//GPU存儲空間int memsize = height * width * sizeof(unsigned char);//輸入 輸出unsigned char* in_gpu;unsigned char* out_gpu;cudaMalloc((void**)&in_gpu, memsize);cudaMalloc((void**)&out_gpu, memsize);cudaEvent_t start, stop_gpu;cudaEventCreate(&start);//創(chuàng)建事件cudaEventCreate(&stop_gpu);//創(chuàng)建事件cudaEventRecord(start);//記錄事件dim3 threadsPreBlock(BLOCK_SIZE, BLOCK_SIZE);dim3 blocksPreGrid((width + threadsPreBlock.x - 1)/threadsPreBlock.x, (height + threadsPreBlock.y - 1)/threadsPreBlock.y);cudaMemcpy(in_gpu, gaussImg.data, memsize, cudaMemcpyHostToDevice);sobel_gpu <<<blocksPreGrid, threadsPreBlock>>> (in_gpu, out_gpu, height, width);cudaEventRecord(stop_gpu);//記錄事件cudaError_t  error_code = cudaGetLastError();if (error_code != cudaSuccess){printf("Error: %s\n", cudaGetErrorString(error_code));printf("FILE: %s\n", __FILE__);printf("LINE: %d\n", __LINE__);printf("Error code: %d\n", error_code);}cudaMemcpy(dst_gpu.data, out_gpu, memsize, cudaMemcpyDeviceToHost);float time_gpu;cudaEventElapsedTime(&time_gpu, start, stop_gpu);//事件計時printf("GPU time: %.3f ms \n", time_gpu);cudaEventDestroy(start);//銷毀事件cudaEventDestroy(stop_gpu);cv::imwrite("dst_gpu_save.png", dst_gpu);//cv::namedWindow("src", cv::WINDOW_NORMAL);cv::imshow("src", src);cv::imshow("dst_gpu", dst_gpu);cv::waitKey();cudaFree(in_gpu);cudaFree(out_gpu);return 0;
}