Mali OpenCL SDK v1.1.0
 All Classes Files Functions Variables Macros Pages
sobel.cpp
Go to the documentation of this file.
1 /*
2  * This confidential and proprietary software may be used only as
3  * authorised by a licensing agreement from ARM Limited
4  * (C) COPYRIGHT 2013 ARM Limited
5  * ALL RIGHTS RESERVED
6  * The entire notice above must be reproduced on all authorised
7  * copies and copies may only be made to the extent permitted
8  * by a licensing agreement from ARM Limited.
9  */
10 
11 #include "common.h"
12 #include "image.h"
13 
14 #include <CL/cl.h>
15 #include <iostream>
16 #include <fstream>
17 #include <sstream>
18 #include <cstddef>
19 #include <cmath>
20 
21 using namespace std;
22 
34 int main(void)
35 {
36  /*
37  * Name of the bitmap to load and run the sobel filter on.
38  * It's width must be divisible by 16.
39  */
40  string filename = "assets/input.bmp";
41 
42  cl_context context = 0;
43  cl_command_queue commandQueue = 0;
44  cl_program program = 0;
45  cl_device_id device = 0;
46  cl_kernel kernel = 0;
47  const unsigned int numberOfMemoryObjects = 3;
48  cl_mem memoryObjects[numberOfMemoryObjects] = {0, 0, 0};
49  cl_int errorNumber;
50 
51  if (!createContext(&context))
52  {
53  cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);
54  cerr << "Failed to create an OpenCL context. " << __FILE__ << ":"<< __LINE__ << endl;
55  return 1;
56  }
57 
58  if (!createCommandQueue(context, &commandQueue, &device))
59  {
60  cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);
61  cerr << "Failed to create the OpenCL command queue. " << __FILE__ << ":"<< __LINE__ << endl;
62  return 1;
63  }
64 
65  if (!createProgram(context, device, "assets/sobel.cl", &program))
66  {
67  cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);
68  cerr << "Failed to create OpenCL program." << __FILE__ << ":"<< __LINE__ << endl;
69  return 1;
70  }
71 
72  kernel = clCreateKernel(program, "sobel", &errorNumber);
73  if (!checkSuccess(errorNumber))
74  {
75  cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);
76  cerr << "Failed to create OpenCL kernel. " << __FILE__ << ":"<< __LINE__ << endl;
77  return 1;
78  }
79 
80  /* Load 24-bits per pixel RGB data from a bitmap. */
81  cl_int width;
82  cl_int height;
83  unsigned char* imageData = NULL;
84  if (!loadFromBitmap(filename, &width, &height, &imageData))
85  {
86  cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);
87  cerr << "Failed loading bitmap. " << __FILE__ << ":"<< __LINE__ << endl;
88  return 1;
89  }
90 
91  /* All buffers are the size of the image data. */
92  size_t bufferSize = width * height * sizeof(cl_uchar);
93 
94  bool createMemoryObjectsSuccess = true;
95  /* Create one input buffer for the image data, and two output buffers for the gradients in X and Y respectively. */
96  memoryObjects[0] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, bufferSize, NULL, &errorNumber);
97  createMemoryObjectsSuccess &= checkSuccess(errorNumber);
98  memoryObjects[1] = clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_ALLOC_HOST_PTR, bufferSize, NULL, &errorNumber);
99  createMemoryObjectsSuccess &= checkSuccess(errorNumber);
100  memoryObjects[2] = clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_ALLOC_HOST_PTR, bufferSize, NULL, &errorNumber);
101  createMemoryObjectsSuccess &= checkSuccess(errorNumber);
102  if (!createMemoryObjectsSuccess)
103  {
104  cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);
105  cerr << "Failed to create OpenCL buffers. " << __FILE__ << ":"<< __LINE__ << endl;
106  return 1;
107  }
108 
109  /* Map the input luminance memory object to a host side pointer. */
110  cl_uchar* luminance = (cl_uchar*)clEnqueueMapBuffer(commandQueue, memoryObjects[0], CL_TRUE, CL_MAP_WRITE, 0, bufferSize, 0, NULL, NULL, &errorNumber);
111  if (!checkSuccess(errorNumber))
112  {
113  cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);
114  cerr << "Mapping memory objects failed " << __FILE__ << ":"<< __LINE__ << endl;
115  return 1;
116  }
117 
118  /* Convert 24-bits per pixel RGB into 8-bits per pixel luminance data. */
119  RGBToLuminance(imageData, luminance, width, height);
120 
121  delete [] imageData;
122 
123  /* Unmap the memory so we can pass it to the kernel. */
124  if (!checkSuccess(clEnqueueUnmapMemObject(commandQueue, memoryObjects[0], luminance, 0, NULL, NULL)))
125  {
126  cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);
127  cerr << "Unmapping memory objects failed " << __FILE__ << ":"<< __LINE__ << endl;
128  return 1;
129  }
130 
131  bool setKernelArgumentsSuccess = true;
132  /* Setup the kernel arguments. */
133  setKernelArgumentsSuccess &= checkSuccess(clSetKernelArg(kernel, 0, sizeof(cl_mem), &memoryObjects[0]));
134  setKernelArgumentsSuccess &= checkSuccess(clSetKernelArg(kernel, 1, sizeof(cl_int), &width));
135  setKernelArgumentsSuccess &= checkSuccess(clSetKernelArg(kernel, 2, sizeof(cl_mem), &memoryObjects[1]));
136  setKernelArgumentsSuccess &= checkSuccess(clSetKernelArg(kernel, 3, sizeof(cl_mem), &memoryObjects[2]));
137  if (!setKernelArgumentsSuccess)
138  {
139  cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);
140  cerr << "Failed setting OpenCL kernel arguments. " << __FILE__ << ":"<< __LINE__ << endl;
141  return 1;
142  }
143 
144  /* An event to associate with the Kernel. Allows us to retreive profiling information later. */
145  cl_event event = 0;
146 
147  /* [Kernel size] */
148  /*
149  * Each instance of the kernel operates on a 16 * 1 portion of the image.
150  * Therefore, the global work size must be width / 16 by height / 1 work items.
151  */
152  size_t globalWorksize[2] = {width / 16, height / 1};
153  /* [Kernel size] */
154 
155  /* Enqueue the kernel */
156  if (!checkSuccess(clEnqueueNDRangeKernel(commandQueue, kernel, 2, NULL, globalWorksize, NULL, 0, NULL, &event)))
157  {
158  cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);
159  cerr << "Failed enqueuing the kernel. " << __FILE__ << ":"<< __LINE__ << endl;
160  return 1;
161  }
162 
163  /* Wait for completion */
164  if (!checkSuccess(clFinish(commandQueue)))
165  {
166  cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);
167  cerr << "Failed waiting for kernel execution to finish. " << __FILE__ << ":"<< __LINE__ << endl;
168  return 1;
169  }
170 
171  /* Print the profiling information for the event. */
172  printProfilingInfo(event);
173  /* Release the event object. */
174  if (!checkSuccess(clReleaseEvent(event)))
175  {
176  cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);
177  cerr << "Failed releasing the event object. " << __FILE__ << ":"<< __LINE__ << endl;
178  return 1;
179  }
180 
181  /* Map the arrays holding the output gradients. */
182  bool mapMemoryObjectsSuccess = true;
183  cl_char* outputDx = (cl_char*)clEnqueueMapBuffer(commandQueue, memoryObjects[1], CL_TRUE, CL_MAP_READ, 0, bufferSize, 0, NULL, NULL, &errorNumber);
184  mapMemoryObjectsSuccess &= checkSuccess(errorNumber);
185  cl_char* outputDy = (cl_char*)clEnqueueMapBuffer(commandQueue, memoryObjects[2], CL_TRUE, CL_MAP_READ, 0, bufferSize, 0, NULL, NULL, &errorNumber);
186  mapMemoryObjectsSuccess &= checkSuccess(errorNumber);
187  if (!mapMemoryObjectsSuccess)
188  {
189  cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);
190  cerr << "Mapping memory objects failed " << __FILE__ << ":"<< __LINE__ << endl;
191  return 1;
192  }
193 
194  /*
195  * In order to better visualise the data, we take the absolute
196  * value of the gradients and then combine them together.
197  * This is work which could be done in the OpenCL kernel, it all depends
198  * on what type of output you require.
199  */
200 
201  /* To visualise the data we take the absolute values of the gradients. */
202  unsigned char *absDX = new unsigned char[width * height];
203  unsigned char *absDY = new unsigned char[width * height];
204  for (int i = 0; i < width * height; i++)
205  {
206  absDX[i] = abs(outputDx[i]);
207  absDY[i] = abs(outputDy[i]);
208  }
209 
210  /* Unmap the memory. */
211  bool unmapMemoryObjectsSuccess = true;
212  unmapMemoryObjectsSuccess &= checkSuccess(clEnqueueUnmapMemObject(commandQueue, memoryObjects[1], outputDx, 0, NULL, NULL));
213  unmapMemoryObjectsSuccess &= checkSuccess(clEnqueueUnmapMemObject(commandQueue, memoryObjects[2], outputDy, 0, NULL, NULL));
214  if (!unmapMemoryObjectsSuccess)
215  {
216  cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);
217  cerr << "Unmapping memory objects failed " << __FILE__ << ":"<< __LINE__ << endl;
218  return 1;
219  }
220 
221  /* Release OpenCL objects. */
222  cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);
223 
224  /* Convert the two output luminance arrays to RGB and save them out to files. */
225  unsigned char* rgbOut = new unsigned char[width * height * 3];
226  luminanceToRGB(absDX, rgbOut, width, height);
227  saveToBitmap("output-dX.bmp", width, height, rgbOut);
228 
229  luminanceToRGB(absDY, rgbOut, width, height);
230  saveToBitmap("output-dY.bmp", width, height, rgbOut);
231 
232  /* Calculate the total gradient of the image, convert it to RGB and store it out to a file. */
233  unsigned char* totalOutput = new unsigned char[width * height];
234  for (int index = 0; index < width * height; index++)
235  {
236  totalOutput[index] = sqrt(pow(absDX[index], 2) + pow(absDY[index], 2));
237  }
238  luminanceToRGB(totalOutput, rgbOut, width, height);
239  saveToBitmap("output.bmp", width, height, rgbOut);
240 
241  delete [] absDX;
242  delete [] absDY;
243  delete [] rgbOut;
244  delete [] totalOutput;
245 
246  return 0;
247 }