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