-
Notifications
You must be signed in to change notification settings - Fork 0
/
Sobel_1.cu
277 lines (239 loc) · 9.28 KB
/
Sobel_1.cu
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
/**************************************************************************
* 文件名:Sobel_1.cu
* 作者: 孙霖(SC19023100)(Seafood)
* 说明: cudau作业之Sobel算子边缘检测第一版
* 调整使用Cpu包、CPU对像素点进行操作、GPU使用1thread、1block对像素点进行操作
****************************************************************************/
//----------------------------头文件包含和空间声明------------------------------------
#include "cuda_runtime.h"
#include "cuda.h"
#include <stdlib.h>
#include <opencv2/opencv.hpp>
#include <iostream>
#include <time.h>
typedef unsigned char uchar;
//命名空间
using namespace std;
using namespace cv;
#define uchar unsigned char
//----------------------------全局变量---------------------------------------------
Mat g_sobelGradient_X, g_sobelGradient_Y; //opencv包使用的全局变量
Mat g_sobelAbsGradient_X, g_sobelAbsGradient_Y; //opencv包使用的全局变量
int g_sobelKernelSize = 1;
Mat dst_Cpu, g_gaussImage, g_grayImage, g_dstImage; //图像读取+高斯滤波
int g_imgHeight, g_imgWidth; //图像的大小Size
//变换矩阵
/**
*@author:Seafood
*@name:sobelInCuda()
*@return:void
*@function:使用Cuda对图像进行Sobel边缘检测
*@para:*dataIn:输入图像 *dataOut:输出图像 imgHeight:图像的高 imgWidth:图像的宽
*其他要注意的地方
**/
//Sobel算子边缘检测核函数
__global__ void sobelInCuda(unsigned char *dataIn, unsigned char *dataOut, int imgHeight, int imgWidth)
{
//用单thread操作
int xIndex = threadIdx.x;
int yIndex = xIndex/imgWidth;
int index = 0;
int Gx = 0;
int Gy = 0;
while(yIndex < imgHeight - 2)
{
Gx = (-1) * dataIn[(yIndex) * imgWidth + xIndex] + 0*dataIn[(yIndex +1 ) * imgWidth + xIndex] + 1*dataIn[(yIndex + 2) * imgWidth + xIndex]
+ (-2) * dataIn[(yIndex) * imgWidth + xIndex + 1] + 0 * dataIn[(yIndex+1) * imgWidth + xIndex + 1] + 2 * dataIn[(yIndex + 2) * imgWidth + xIndex + 1]
+ (-1) * dataIn[(yIndex) * imgWidth + xIndex + 2] + 0 * dataIn[(yIndex+1) * imgWidth + xIndex + 2] + 1 * dataIn[(yIndex + 2) * imgWidth + xIndex + 2];
Gy = (-1) * dataIn[(yIndex) * imgWidth + xIndex] + (-2) * dataIn[(yIndex +1 ) * imgWidth + xIndex] + (-1)*dataIn[(yIndex + 2) * imgWidth + xIndex]
+ (0) * dataIn[(yIndex) * imgWidth + xIndex + 1] + 0 * dataIn[(yIndex+1) * imgWidth + xIndex + 1] + 0 * dataIn[(yIndex + 2) * imgWidth + xIndex + 1]
+ (1) * dataIn[(yIndex) * imgWidth + xIndex + 2] + 2 * dataIn[(yIndex+1) * imgWidth + xIndex + 2] + 1 * dataIn[(yIndex + 2) * imgWidth + xIndex + 2];
int sum = abs(Gx) + abs(Gy);
if (sum > 255)
{
sum = 255; //for best performance
}
dataOut[index] = sum;
xIndex ++;
if ( xIndex == imgWidth-2 )
{
xIndex = 0;
yIndex ++;
}
index = xIndex + yIndex * imgWidth;
}
}
/**
*@author:Seafood
*@name:sobelCpuPixel()
*@return:void
*@function:使用CPU对图像像素进行Sobel边缘检测
*@para:None
*其他要注意的地方
**/
//Sobel算子边缘检测CPU函数
void sobelCpuPixel()
{
clock_t begin_time, end_clock;
begin_time = clock();
int sobel_x[3][3];
int sobel_y[3][3];
// asigning values to sobel x direction
sobel_x[0][0] = -1; sobel_x[0][1] = 0; sobel_x[0][2] =1;
sobel_x[1][0] = -2; sobel_x[1][1] = 0; sobel_x[1][2] =2;
sobel_x[2][0] = -1; sobel_x[2][1] = 0; sobel_x[2][2] =1;
// asigning values to sobel y direction
sobel_y[0][0] = -1; sobel_y[0][1] = -2; sobel_y[0][2] = -1;
sobel_y[1][0] = 0; sobel_y[1][1] = 0; sobel_y[1][2] = 0;
sobel_y[2][0] = 1; sobel_y[2][1] = 2; sobel_y[2][2] = 1;
Mat img = g_gaussImage;
Mat newimg = img;
for (int j = 0; j<img.rows-2; j++)
{
for (int i = 0; i<img.cols-2; i++)
{
int pixval_x =
(sobel_x[0][0] * (int)img.at<uchar>(j,i)) + (sobel_x[0][1] * (int)img.at<uchar>(j+1,i)) + (sobel_x[0][2] * (int)img.at<uchar>(j+2,i)) +
(sobel_x[1][0] * (int)img.at<uchar>(j,i+1)) + (sobel_x[1][1] * (int)img.at<uchar>(j+1,i+1)) + (sobel_x[1][2] * (int)img.at<uchar>(j+2,i+1)) +
(sobel_x[2][0] * (int)img.at<uchar>(j,i+2)) + (sobel_x[2][1] * (int)img.at<uchar>(j+1,i+2)) + (sobel_x[2][2] * (int)img.at<uchar>(j+2,i+2));
int pixval_y =
(sobel_y[0][0] * (int)img.at<uchar>(j,i)) + (sobel_y[0][1] * (int)img.at<uchar>(j+1,i)) + (sobel_y[0][2] * (int)img.at<uchar>(j+2,i)) +
(sobel_y[1][0] * (int)img.at<uchar>(j,i+1)) + (sobel_y[1][1] * (int)img.at<uchar>(j+1,i+1)) + (sobel_y[1][2] * (int)img.at<uchar>(j+2,i+1)) +
(sobel_y[2][0] * (int)img.at<uchar>(j,i+2)) + (sobel_y[2][1] * (int)img.at<uchar>(j+1,i+2)) + (sobel_y[2][2] * (int)img.at<uchar>(j+2,i+2));
int sum = abs(pixval_x) + abs(pixval_y);
if (sum > 255)
{
sum = 255; //for best performance
}
newimg.at<uchar>(j,i) = sum;
}
}
end_clock = clock();
cout << "CPU对像素操作运行时间为: " << static_cast<double>(end_clock - begin_time) / CLOCKS_PER_SEC*1000 << "ms" << endl;//输出运行时间为毫秒
cvNamedWindow("processed_by_CPUinPaxel", 0);
resizeWindow("processed_by_CPUinPaxel", 800, 600);
imshow("processed_by_CPUinPaxel", newimg);
}
/**
*@author:Seafood
*@name:sobelCpuPackage
*@return:void
*@function:使用CPU的OPENCV包对图像进行Sobel边缘检测
*@para:None
*其他要注意的地方
**/
//Sobel算子边缘检测CPU函数
void sobelCpuPackage(int, void*)
{
clock_t begin_time, end_clock;
begin_time = clock();
//求x方向梯度
Sobel(g_grayImage, g_sobelGradient_X, CV_16S, 1, 0, (2*g_sobelKernelSize + 1), 1, 1, BORDER_DEFAULT);
convertScaleAbs( g_sobelGradient_X, g_sobelAbsGradient_X);
//求Y方向梯度
Sobel(g_grayImage, g_sobelGradient_Y, CV_16S, 1, 0, (2*g_sobelKernelSize + 1), 1, 1, BORDER_DEFAULT);
convertScaleAbs( g_sobelGradient_Y, g_sobelAbsGradient_Y);
//合并梯度
addWeighted(g_sobelAbsGradient_X, 0.5, g_sobelAbsGradient_Y, 0.5,0,dst_Cpu );
end_clock = clock();
cout << "CPU对使用OPENCV包操作运行时间为: " << static_cast<double>(end_clock - begin_time) / CLOCKS_PER_SEC*1000 << "ms" << endl;//输出运行时间为毫秒
cvNamedWindow("CPU处理效果图(OPENCV包)", 0);
resizeWindow("CPU处理效果图(OPENCV包)", 800, 600);
imshow("CPU处理效果图(OPENCV包)", dst_Cpu);
}
/**
*@author:Seafood
*@name:pictureInit(void)
*@return:void
*@function:图像读取、高斯滤波和初始化
*@para:None
*其他要注意的地方
**/
//图像读入
void pictureInit(void)
{
//读入图像
g_grayImage = imread("ironman.jpg", 0);
//显示原图
cvNamedWindow("originimage", 0);
resizeWindow("originimage", 800, 600);
imshow("originimage", g_grayImage);
g_imgHeight = g_grayImage.rows;
g_imgWidth = g_grayImage.cols;
printf(" picture size is %d x %d \n", g_imgHeight, g_imgWidth);
//高斯滤波
GaussianBlur(g_grayImage, g_gaussImage, Size(3,3), 0, 0, BORDER_DEFAULT);
g_dstImage = g_gaussImage;
}
/**
*@author:Seafood
*@name:sobelGPUPixel()
*@return:void
*@function:调用GPU对图像处理并计算时间
*@para:None
*其他要注意的地方
**/
//GPU处理函数
void sobelGPUPixel()
{
unsigned char *dev_in;
unsigned char *dev_out;
//创建时间用于计算
cudaEvent_t start, stop;
cudaEventCreate( &start );
cudaEventCreate( &stop );
//开始时间点记录
cudaEventRecord( start, 0 );
//申请内存
cudaMalloc((void**)&dev_in, g_imgHeight * g_imgWidth * sizeof(unsigned char));
cudaMalloc((void**)&dev_out, g_imgHeight * g_imgWidth * sizeof(unsigned char));
Mat dst_Gpu(g_imgHeight, g_imgWidth, CV_8UC1, Scalar(0));
//导入内存
cudaMemcpy(dev_in, g_gaussImage.data, g_imgHeight * g_imgWidth * sizeof(unsigned char), cudaMemcpyHostToDevice);
//调用核函数
//单block 单thread
sobelInCuda<< <1,1>> >(dev_in, dev_out, g_imgHeight, g_imgWidth);
//导出处理
cudaMemcpy(dst_Gpu.data, dev_out, g_imgHeight * g_imgWidth * sizeof(unsigned char), cudaMemcpyDeviceToHost);
//停止时间点
cudaEventRecord( stop, 0 );
cudaEventSynchronize( stop );
//释放内存
cudaFree(dev_in);
cudaFree(dev_out);
//计算GPU所用时间
float elapsedTime;
cudaEventElapsedTime( &elapsedTime, start, stop );
printf( "GPU对像素操作运行时间为: %.1f ms \n", elapsedTime );
cudaEventDestroy( start );
cudaEventDestroy( stop );
//显示处理后的图像
cvNamedWindow("GPU处理后的图像", 0);
resizeWindow("GPU处理后的图像", 800, 600);
imshow("GPU处理后的图像", dst_Gpu);
}
/**
*@author:Seafood
*@name:main(int argc, char *argv[], char **env)
*@return:int
*@function:程序入口,主函数
*@para:None
*其他要注意的地方
**/
//main函数
int main(int argc, char *argv[], char **env)
{
//图像读入和处理
pictureInit();
//Sobel算子Cpu package实现
sobelCpuPackage(0,0);
//Sobel算子Cpu 对像素操作实现
sobelCpuPixel();
//Sobel算子GPU操作实现
sobelGPUPixel();
//结束展示
while((char)waitKey(0) != 'q' )
{
}
return 0;
}