Mali OpenCL SDK v1.1.0
 All Classes Files Functions Variables Macros Pages
image_scaling.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 
17 using namespace std;
18 
24 int main(void)
25 {
26  cl_context context = 0;
27  cl_command_queue commandQueue = 0;
28  cl_program program = 0;
29  cl_device_id device = 0;
30  cl_kernel kernel = 0;
31  const int numMemoryObjects = 2;
32  cl_mem memoryObjects[numMemoryObjects] = {0, 0};
33  cl_int errorNumber;
34 
35  /* Set up OpenCL environment: create context, command queue, program and kernel. */
36  if (!createContext(&context))
37  {
38  cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numMemoryObjects);
39  cerr << "Failed to create an OpenCL context. " << __FILE__ << ":"<< __LINE__ << endl;
40  return 1;
41  }
42 
43  if (!createCommandQueue(context, &commandQueue, &device))
44  {
45  cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numMemoryObjects);
46  cerr << "Failed to create the OpenCL command queue. " << __FILE__ << ":"<< __LINE__ << endl;
47  return 1;
48  }
49 
50  if (!createProgram(context, device, "assets/image_scaling.cl", &program))
51  {
52  cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numMemoryObjects);
53  cerr << "Failed to create OpenCL program." << __FILE__ << ":"<< __LINE__ << endl;
54  return 1;
55  }
56 
57  kernel = clCreateKernel(program, "image_scaling", &errorNumber);
58  if (!checkSuccess(errorNumber))
59  {
60  cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numMemoryObjects);
61  cerr << "Failed to create OpenCL kernel. " << __FILE__ << ":"<< __LINE__ << endl;
62  return 1;
63  }
64 
65  /* Print the image formats that the OpenCL device supports. */
66  cout << endl;
68  cout << endl;
69 
70  /* The scaling factor to use when resizing the image. */
71  const int scaleFactor = 8;
72 
73  /* Load the input image data. */
74  unsigned char* inputImage = NULL;
75  int width, height;
76  loadFromBitmap("assets/input.bmp", &width, &height, &inputImage);
77 
78  /*
79  * Calculate the width and height of the new image.
80  * Used to allocate the correct amount of output memory and the number of kernels to use.
81  */
82  int newWidth = width * scaleFactor;
83  int newHeight = height * scaleFactor;
84 
85  /* [Allocate image objects] */
86  /*
87  * Specify the format of the image.
88  * The bitmap image we are using is RGB888, which is not a supported OpenCL image format.
89  * We will use RGBA8888 and add an empty alpha channel.
90  */
91  cl_image_format format;
92  format.image_channel_data_type = CL_UNORM_INT8;
93  format.image_channel_order = CL_RGBA;
94 
95  /* Allocate memory for the input image that can be accessed by the CPU and GPU. */
96  bool createMemoryObjectsSuccess = true;
97 
98  memoryObjects[0] = clCreateImage2D(context, CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, &format, width, height, 0, NULL, &errorNumber);
99  createMemoryObjectsSuccess &= checkSuccess(errorNumber);
100 
101  memoryObjects[1] = clCreateImage2D(context, CL_MEM_WRITE_ONLY | CL_MEM_ALLOC_HOST_PTR, &format, newWidth, newHeight, 0, NULL, &errorNumber);
102  createMemoryObjectsSuccess &= checkSuccess(errorNumber);
103 
104  if (!createMemoryObjectsSuccess)
105  {
106  cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numMemoryObjects);
107  cerr << "Failed creating the image. " << __FILE__ << ":"<< __LINE__ << endl;
108  return 1;
109  }
110  /* [Allocate image objects] */
111 
112  /* [Map image objects to host pointers] */
113  /*
114  * Like with memory buffers, we now map the allocated memory to a host side pointer.
115  * Unlike buffers, we must specify origin coordinates, width and height for the region of the image we wish to map.
116  */
117  size_t origin[3] = {0, 0, 0};
118  size_t region[3] = {width, height, 1};
119 
120  /*
121  * clEnqueueMapImage also returns the rowPitch; the width of the mapped region in bytes.
122  * If the image format is not known, this is required information when accessing the image object as a normal array.
123  * The number of bytes per pixel can vary with the image format being used,
124  * this affects the offset into the array for a given coordinate.
125  * In our case the image format is fixed as RGBA8888 so we don't need to worry about the rowPitch.
126  */
127  size_t rowPitch;
128 
129  unsigned char* inputImageRGBA = (unsigned char*)clEnqueueMapImage(commandQueue, memoryObjects[0], CL_TRUE, CL_MAP_WRITE, origin, region, &rowPitch, NULL, 0, NULL, NULL, &errorNumber);
130  if (!checkSuccess(errorNumber))
131  {
132  cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numMemoryObjects);
133  cerr << "Failed mapping the input image. " << __FILE__ << ":"<< __LINE__ << endl;
134  return 1;
135  }
136  /* [Map image objects to host pointers] */
137 
138  /* Convert the input data from RGB to RGBA (moves it to the OpenCL allocated memory at the same time). */
139  RGBToRGBA(inputImage, inputImageRGBA, width, height);
140  delete[] inputImage;
141 
142  /* Unmap the image from the host. */
143  if (!checkSuccess(clEnqueueUnmapMemObject(commandQueue, memoryObjects[0], inputImageRGBA, 0, NULL, NULL)))
144  {
145  cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numMemoryObjects);
146  cerr << "Failed unmapping the input image. " << __FILE__ << ":"<< __LINE__ << endl;
147  return 1;
148  }
149 
150  /*
151  * Calculate the normalization factor for the image coordinates.
152  * By using normalized coordinates we don't have to manually map the destination coordinates to the source coordinates.
153  */
154  cl_float widthNormalizationFactor = 1.0f / newWidth;
155  cl_float heightNormalizationFactor = 1.0f / newHeight;
156 
157  /* Setup the kernel arguments. */
158  bool setKernelArgumentsSuccess = true;
159  setKernelArgumentsSuccess &= checkSuccess(clSetKernelArg(kernel, 0, sizeof(cl_mem), &memoryObjects[0]));
160  setKernelArgumentsSuccess &= checkSuccess(clSetKernelArg(kernel, 1, sizeof(cl_mem), &memoryObjects[1]));
161  setKernelArgumentsSuccess &= checkSuccess(clSetKernelArg(kernel, 2, sizeof(cl_float), &widthNormalizationFactor));
162  setKernelArgumentsSuccess &= checkSuccess(clSetKernelArg(kernel, 3, sizeof(cl_float), &heightNormalizationFactor));
163  if (!setKernelArgumentsSuccess)
164  {
165  cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, 3);
166  cerr << "Failed setting OpenCL kernel arguments. " << __FILE__ << ":"<< __LINE__ << endl;
167  return 1;
168  }
169 
170  /*
171  * Set the kernel work size. Each kernel operates on one pixel of the ouput image.
172  * Therefore, we need newWidth * newHeight kernel instances.
173  * We are using two work dimensions because it maps nicely onto the coordinates of the image.
174  * With one dimension we would have to derive the y coordinate from the x coordinate in the kernel.
175  */
176  const int workDimensions = 2;
177  size_t globalWorkSize[workDimensions] = {newWidth, newHeight};
178 
179  /* An event to associate with the kernel. Allows us to retrieve profiling information later. */
180  cl_event event = 0;
181 
182  /* Enqueue the kernel. */
183  if (!checkSuccess(clEnqueueNDRangeKernel(commandQueue, kernel, workDimensions, NULL, globalWorkSize, NULL, 0, NULL, &event)))
184  {
185  cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numMemoryObjects);
186  cerr << "Failed enqueuing the kernel. " << __FILE__ << ":"<< __LINE__ << endl;
187  return 1;
188  }
189 
190  /* Wait for kernel execution completion. */
191  if (!checkSuccess(clFinish(commandQueue)))
192  {
193  cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numMemoryObjects);
194  cerr << "Failed waiting for kernel execution to finish. " << __FILE__ << ":"<< __LINE__ << endl;
195  return 1;
196  }
197 
198  /* Print the profiling information for the event. */
199  printProfilingInfo(event);
200  /* Release the event object. */
201  if (!checkSuccess(clReleaseEvent(event)))
202  {
203  cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numMemoryObjects);
204  cerr << "Failed releasing the event object. " << __FILE__ << ":"<< __LINE__ << endl;
205  return 1;
206  }
207 
208  size_t newRegion[3] = {newWidth, newHeight, 1};
209 
210  unsigned char* outputImage = (unsigned char*)clEnqueueMapImage(commandQueue, memoryObjects[1], CL_TRUE, CL_MAP_READ, origin, newRegion, &rowPitch, NULL, 0, NULL, NULL, &errorNumber);
211  if (!checkSuccess(errorNumber))
212  {
213  cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numMemoryObjects);
214  cerr << "Failed mapping the input image. " << __FILE__ << ":"<< __LINE__ << endl;
215  return 1;
216  }
217 
218  unsigned char* outputImageRGB = new unsigned char[newWidth * newHeight * 3];
219  RGBAToRGB(outputImage, outputImageRGB, newWidth, newHeight);
220 
221  saveToBitmap("output.bmp", newWidth, newHeight, outputImageRGB);
222 
223  delete[] outputImageRGB;
224 
225  cleanUpOpenCL(context, commandQueue, program, kernel, memoryObjects, numMemoryObjects);
226 
227 
228  return 0;
229 }