4、host端程序代码
Host端程序处理流程就是按照前面“程序设计”一节编写的。除了调用OpenCL+OpenCV的API函数,其他的地方都是按照C/C++语法编写的。
具体代码如下:
1. // ImageRotate.cpp : 定义控制台应用程序的入口点。 2. // 3. 4. #include "stdafx.h" 5. #include <iostream> 6. #include <fstream> 7. #include <sstream> 8. 9. #include <opencv2/opencv.hpp> 10. 11. #ifdef __APPLE__ 12. #include <OpenCL/cl.h> 13. #else 14. #include <CL/cl.h> 15. #endif 16. 17. using namespace cv; 18. 19. int _tmain(int argc, _TCHAR* argv[]) 20. { 21. cl_int ciErrNum; 22. const char *fileName = "ImageRotate.cl"; 23. int width = 0, height = 0; 24. float cos_theta = 0.7071067811865476, sin_theta = 0.7071067811865475; //for degree 45 25. //float cos_theta = 0.5, sin_theta = 0.5; 26. const char* imageName = "F:\\code\\pic\\test01.jpg"; 27. char *bufInput = NULL, *bufOutput = NULL; 28. 29. 30. //read one jpeg pic and store it in a Mat variable. 31. Mat img = imread(imageName); 32. if (!img.data) { 33. std::cout << "fail to open the file:" << imageName << std::endl; 34. } 35. 36. //the type of img is RGB, convert to gray image. 37. Mat imgGray; 38. cvtColor(img, imgGray, CV_BGR2GRAY); 39. width = imgGray.cols; 40. height = imgGray.rows; 41. std::cout << "picture width: " << width << ", height: " << height << std::endl; 42. 43. //save the source data of original gray image. 44. FILE *yuvFileOrg = NULL; 45. fopen_s(&yuvFileOrg, "gray_org.yuv", "wb"); 46. fwrite(imgGray.data, width * height * sizeof(unsigned char), 1, yuvFileOrg); 47. fclose(yuvFileOrg); 48. yuvFileOrg = NULL; 49. 50. //display the original gray image in a window. 51. namedWindow( imageName, CV_WINDOW_AUTOSIZE ); 52. imshow(imageName, imgGray); 53. //waitKey(0); 54. 55. //allocate the input buffer to store the original gray image 56. if (NULL == (bufInput = (char *)malloc(width * height * sizeof(char)))) { 57. std::cerr << "Failed to malloc buffer for input image. " << std::endl; 58. return NULL; 59. } 60. 61. //allocate the output buffer to store the image rotated. 62. if (NULL == (bufOutput = (char *)malloc(width * height * sizeof(char)))) { 63. std::cerr << "Failed to malloc buffer for output image. " << std::endl; 64. return NULL; 65. } 66. 67. //copy the data of gray image to the input buffer. initialize the output buffer by zero. 68. memcpy(bufInput, imgGray.data, width * height * sizeof(unsigned char)); 69. memset(bufOutput, 0, width * height * sizeof(unsigned char)); 70. 71. //use the first platform 72. cl_platform_id platform; 73. ciErrNum = clGetPlatformIDs(1, &platform, NULL); 74. 75. //use the first device 76. cl_device_id device; 77. ciErrNum = clGetDeviceIDs( 78. platform, 79. CL_DEVICE_TYPE_ALL, 80. 1, 81. &device, 82. NULL); 83. 84. cl_context_properties cps[3] = { 85. CL_CONTEXT_PLATFORM, (cl_context_properties)platform, 0 86. }; 87. //create the context 88. cl_context ctx = clCreateContext( 89. cps, 90. 1, 91. &device, 92. NULL, 93. NULL, 94. &ciErrNum); 95. 96. //create the command queue 97. cl_command_queue myqueue = clCreateCommandQueue( 98. ctx, 99. device, 100. 0, 101. &ciErrNum); 102. 103. //allocate space for original image on the device 104. cl_mem bufferA = clCreateBuffer( 105. ctx, 106. CL_MEM_READ_ONLY, 107. width * height * sizeof(unsigned char), 108. NULL, 109. &ciErrNum); 110. //copy input buffer to the device 111. ciErrNum = clEnqueueWriteBuffer( 112. myqueue, 113. bufferA, 114. CL_TRUE, 115. 0, 116. width * height * sizeof(unsigned char), 117. (void *)bufInput, 118. 0, 119. NULL, 120. NULL); 121. 122. //allocate space for rotated image on the device 123. cl_mem bufferC = clCreateBuffer( 124. ctx, 125. CL_MEM_READ_WRITE, 126. width * height * sizeof(unsigned char), 127. NULL, 128. &ciErrNum); 129. 130. //open kernel file and read the content to a string variable. 131. std::ifstream kernelFile("ImageRotate.cl", std::ios::in); 132. if (!kernelFile.is_open()) { 133. std::cerr << "Failed to open file for reading: " << fileName << std::endl; 134. return NULL; 135. } 136. std::ostringstream oss; 137. oss << kernelFile.rdbuf(); 138. std::string srcStdStr = oss.str(); 139. const char *srcStr = srcStdStr.c_str(); 140. kernelFile.close(); 141. 142. //create the program with source code of kernel. 143. cl_program myprog = clCreateProgramWithSource( 144. ctx, 145. 1, 146. (const char**)&srcStr, 147. NULL, 148. &ciErrNum); 149. 150. //compile the program. passing NULL for the 'device_list' argument targets all devices in the context 151. ciErrNum = clBuildProgram(myprog, 0, NULL, NULL, NULL, NULL); 152. 153. //create the kernel 154. cl_kernel mykernel = clCreateKernel( 155. myprog, 156. "img_rotate", 157. &ciErrNum); 158. 159. //set the kernel arguments 160. clSetKernelArg(mykernel, 0, sizeof(cl_mem), (void *)&bufferC); 161. clSetKernelArg(mykernel, 1, sizeof(cl_mem), (void *)&bufferA); 162. clSetKernelArg(mykernel, 2, sizeof(cl_int), (void *)&width); 163. clSetKernelArg(mykernel, 3, sizeof(cl_int), (void *)&height); 164. clSetKernelArg(mykernel, 4, sizeof(cl_float), (void *)&cos_theta); 165. clSetKernelArg(mykernel, 5, sizeof(cl_float), (void *)&sin_theta); 166. 167. //set local and global workgroup sizes 168. size_t localws[2] = {1, 1}; 169. size_t globalws[2] = {width, height}; 170. 171. //execute the kernel 172. ciErrNum = clEnqueueNDRangeKernel( 173. myqueue, 174. mykernel, 175. 2, 176. NULL, 177. globalws, 178. localws, 179. 0, 180. NULL, 181. NULL); 182. 183. //read the output data back to the host 184. ciErrNum = clEnqueueReadBuffer( 185. myqueue, 186. bufferC, 187. CL_TRUE, 188. 0, 189. width * height * sizeof(unsigned char), 190. bufOutput, 191. 0, 192. NULL, 193. NULL); 194. 195. //copy the output data from output buffer to Mat variable. 196. memcpy(imgGray.data, bufOutput, width * height * sizeof(unsigned char)); 197. 198. //save the source data for gray image rotated 199. FILE *yuvFile = NULL; 200. fopen_s(&yuvFile, "gray.yuv", "wb"); 201. fwrite(imgGray.data, width * height * sizeof(unsigned char), 1, yuvFile); 202. fclose(yuvFile); 203. yuvFile = NULL; 204. 205. //save the gray image rotated. 206. imwrite("test_gray.jpg", imgGray); 207. 208. //show the gray image rotated. 209. const char *winName = "gray_image_rotated"; 210. namedWindow(winName, CV_WINDOW_AUTOSIZE ); 211. imshow(winName, imgGray); 212. waitKey(0); 213. destroyAllWindows(); 214. 215. //release all resource 216. if (bufInput != NULL) 217. free(bufInput); 218. 219. if (bufOutput != NULL) 220. free(bufOutput); 221. 222. if (bufferA != 0) 223. clReleaseMemObject(bufferA); 224. 225. if (bufferC != 0) 226. clReleaseMemObject(bufferC); 227. 228. if (myqueue != 0) 229. clReleaseCommandQueue(myqueue); 230. 231. if (mykernel != 0) 232. clReleaseKernel(mykernel); 233. 234. if (myprog != 0) 235. clReleaseProgram(myprog); 236. 237. if (ctx != 0) 238. clReleaseContext(ctx); 239. 240. return 0; 241. }
5、程序处理结果
原始的灰度图像如下所示:
经过45度旋转的图像如下:
将sin、cos值都设置为0.5时的处理结果如下:
经过45度旋转的图像上面有很多暗点,那些暗点是在设备端分配的buffer中原始的点。意味着那些区域并非每个点都有对应的点旋转过来。
在kernel程序中计算坐标点时,使用的是float类型,最后获取图像时是将float转换为int了。应该是因为精度损失导致某些坐标点上没有对应的像素点,所以保留了buffer中原有的数据。
而当sin、cos设置为0.5时,图像看着就没有暗纹。按照0.5计算很少有精度的损失,所以就不存在上面的问题。
是因为旋转算法精度不够,还是程序哪里实现错了,这是个问题,将在后面的学习中去寻找问题的答案。
(完)