Mali OpenCL SDK v1.1.0 教程样例之六“索贝尔滤波器”

Sobel图像滤波器是一种简单的卷积滤波器，主要用于边缘检测算法。

样例结果

输入图像

输出图像

算法

　　一种做图像边缘检测的技术是，找出图像的梯度。大梯度值的区域对应图像色彩或密度的剧变区域。典型地，这些区域是边缘。

　　如果你对于幅图像卷积两个Sobel算子，你会得到两个输出：

　> X方向的梯度(dX)

　> Y方向的梯度(dY)

 

　　通常梯度被合并，给出一个总的梯度图。

　　Sobel算子：dX与dY：

[cpp]  view plain copy print ?

1. -1  0  1  
2. -2  0  2  
3. -1  0  1  

[cpp]  view plain copy print ?

1.  1  2  1  
2.  0  0  0  
3. -1 -2 -1  

 

 

实现

填充

　　值得注意的是，我们在此并不考虑填充。在两个维度，输出图像相比小两个像素。因为，每个输出需要围绕它的像素，计算边沿像素的输出是不可能的。在此样例中，我们简单地把边缘输出像素用其初始值。

　　有时需要输出信号的大小与输入信号的大小一样，这种种情况下，"填充"必须应用到输入中，考虑到滤波应用的实际本性上是减少大小的。填充策略会有不同，但是对于图像，通常的选择是在所有边上重复边缘值(例如，最外层的像素集)，或者在一些边上(在某些情况下)。

实现细节

　　简单起见，我们在一个8比特通道上实现Sobel滤波器。为了在RGB图像上做Sobel滤波，你可以独立地在每个通道上运行Sobel滤波，然后组合结果。在这个例子中，我们取了一张RGB图像，转换成8位辉度图，然后发送到GPU。

　　每个Sobel计算给出一个中心掩码像素的输出。中心像素的输出值是围绕像素的3x3网格上的像素值乘以Sobel掩码的和。这可以被分成三个阶段，通过为网格上的每一行独立地做乘法和累加(这一方法被使用在这个样例中)。

代码

　　注意：我们考虑矢量化是在本教程中使用的主要优化技术。为了显示使用适量的重要性，一个不包含矢量的Sobel滤波器的OpenCL实现被包含在sobel_no_vectors sample(sobel_no_vectors.cpp和sobel_no_vectors.cl)。在Mali-T600系列GPU上，矢量化的版本快很多。

　　除非另作说明，否则所有的代码片段来自sobel.cl中的OpenCL内核。

1. 选择内核大小

　　我们在内核中使用向量类型，因此我们实际上每个内核输出16个结果。看下面更多的向量化细节。内核使用3x3的Sobel滤波器到一个输入图像中的18x3的窗，在两个输出图像中来生成16x1结果，代表梯度的dx和dy。

[cpp]  view plain copy print ?

1. /* 
2.  * Each kernel calculates 16 output pixels in the same row (hence the '* 16'). 
3.  * column is in the range [0, width] in steps of 16. 
4.  * row is in the range [0, height]. 
5.  */  
6. const int column = get_global_id(0) * 16;  
7. const int row = get_global_id(1) * 1;  
8. /* Offset calculates the position in the linear data for the row and the column. */  
9. const int offset = row * width + column;  

　 　当我们在sobel.cpp中排队内核时，相应地减少工作大小：

[cpp]  view plain copy print ?

1. /* 
2.  * Each instance of the kernel operates on a 16 * 1 portion of the image. 
3.  * Therefore, the global work size must be width / 16 by height / 1 work items. 
4.  */  
5. size_t globalWorksize[2] = {width / 16, height / 1};  

2. 加载输入数据

　　Mali-T600系列GPU具有128位矢量寄存器，可以在矢量类型上做运算。因此，我们使用OpenCL矢量，来使得更加有效地使用硬件，从而有更高的性能。

　　这里我们从一行数据中做矢量加载：

[cpp]  view plain copy print ?

1. /* 
2.  * First row of input. 
3.  * In a scalar Sobel calculation you would load 1 value for leftLoad, middleLoad and rightLoad. 
4.  * In the vector case we load 16 values for each. 
5.  * leftLoad, middleLoad and rightLoad load 16-bytes of data from the first row. 
6.  * The data they load overlaps. e.g. for the first column and row, leftLoad is 0->15, middleLoad is 1->16 and rightLoad is 2->17. 
7.  * So we're actually loading 18-bytes of data from the first row. 
8.  */  
9. uchar16 leftLoad = vload16(0, inputImage + (offset + 0));  
10. uchar16 middleLoad = vload16(0, inputImage + (offset + 1));  
11. uchar16 rightLoad = vload16(0, inputImage + (offset + 2));  

3. 转换数据

　　在Mali-T600系列GPU上，扩展和收缩数据类型是一个很容易地操作。在此，我们将8比特每像素数据转换成16比特每像素。

[cpp]  view plain copy print ?

1. /* 
2.  * Convert the data from unsigned chars to shorts (8-bit unsigned to 16-bit signed). 
3.  * The calculations can overflow 8-bits so we require larger intermediate storage. 
4.  * Additionally, the values can become negative so we need a signed type. 
5.  */  
6. short16 leftData = convert_short16(leftLoad);  
7. short16 middleData = convert_short16(middleLoad);  
8. short16 rightData = convert_short16(rightLoad);  

4. 做计算

　　然后，我们在计算16个像素的结果。在Mali-T600系列GPU上，每个向量计算可以作为一个单一操作完成。

[cpp]  view plain copy print ?

1. /* 
2.  * Calculate the results for the first row. 
3.  * Looking at the Sobel masks above for the first line of input, 
4.  * the dX calculation is the sum of 1 * leftData, 0 * middleData, and -1 * rightData. 
5.  * The dY calculation is the sum of 1 * leftData, 2 * middleData, and 1 * rightData. 
6.  * This is what is being calculated below, except we have removed the 
7.  * unnecessary calculations (multiplications by 1 or 0) and we are calculating 16 values at once. 
8.  * This pattern repeats for the other 2 rows of data. 
9.  */  
10. short16 dx = rightData - leftData;  
11. short16 dy = rightData + leftData + middleData * (short)2;  

5. 排列结果

　　最后，我们收缩和存储数据。我们使用一个向量存储一次写16个结果：

[cpp]  view plain copy print ?

1. /* 
2.  * Store the results. 
3.  * The range of outputs from our Sobel calculations is [-1020, 1020]. 
4.  * In order to output this as an 8-bit signed char we must divide it by 8 (or shift right 3 times). 
5.  * This gives the range [-128, 128]. Depending on what type of output you require, 
6.  * (signed/unsigned, seperate/combined gradients) it is possible to do more of the calculations on the GPU using OpenCL. 
7.  * In this sample we're assuming that the application requires signed uncombined gradient outputs. 
8.  */  
9. vstore16(convert_char16(dx >> 3), 0, outputImageDX + offset + width + 1);  
10. vstore16(convert_char16(dy >> 3), 0, outputImageDY + offset + width + 1);  

　　因为数据作为两个分开的梯度图返回，我们在CPU上将它们先组合，然后将它们作为一个位图写入。

运行样例

　　参见样例1。

 

 

附录1：内核源码

[cpp]  view plain copy print ?

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. /** 
12.  * \brief Sobel filter kernel function. 
13.  * \param[in] inputImage Input image data in row-major format. 
14.  * \param[in] width Width of the image passed in as inputImage. 
15.  * \param[out] outputImageDX Output image of the calculated gradient in the X direction. 
16.  * \param[out] outputImageDY Output image of the calculated gradient in the Y direction. 
17.  */  
18. __kernel void sobel(__global const uchar* restrict inputImage,  
19.                     const int width,  
20.                     __global char* restrict outputImageDX,  
21.                     __global char* restrict outputImageDY)  
22. {  
23.     /* [Kernel size] */  
24.     /* 
25.      * Each kernel calculates 16 output pixels in the same row (hence the '* 16'). 
26.      * column is in the range [0, width] in steps of 16. 
27.      * row is in the range [0, height]. 
28.      */  
29.     const int column = get_global_id(0) * 16;  
30.     const int row = get_global_id(1) * 1;  
31.   
32.     /* Offset calculates the position in the linear data for the row and the column. */  
33.     const int offset = row * width + column;  
34.     /* [Kernel size] */  
35.   
36.     /* [Load row] */  
37.     /* 
38.      * First row of input. 
39.      * In a scalar Sobel calculation you would load 1 value for leftLoad, middleLoad and rightLoad. 
40.      * In the vector case we load 16 values for each. 
41.      * leftLoad, middleLoad and rightLoad load 16-bytes of data from the first row. 
42.      * The data they load overlaps. e.g. for the first column and row, leftLoad is 0->15, middleLoad is 1->16 and rightLoad is 2->17. 
43.      * So we're actually loading 18-bytes of data from the first row. 
44.      */  
45.     uchar16 leftLoad = vload16(0, inputImage + (offset + 0));  
46.     uchar16 middleLoad = vload16(0, inputImage + (offset + 1));  
47.     uchar16 rightLoad = vload16(0, inputImage + (offset + 2));  
48.     /* [Load row] */  
49.   
50.     /* [Convert data] */  
51.     /* 
52.      * Convert the data from unsigned chars to shorts (8-bit unsigned to 16-bit signed). 
53.      * The calculations can overflow 8-bits so we require larger intermediate storage. 
54.      * Additionally, the values can become negative so we need a signed type. 
55.      */  
56.     short16 leftData = convert_short16(leftLoad);  
57.     short16 middleData = convert_short16(middleLoad);  
58.     short16 rightData = convert_short16(rightLoad);  
59.     /* [Convert data] */  
60.   
61.     /* [Calculation] */  
62.     /* 
63.      * Calculate the results for the first row. 
64.      * Looking at the Sobel masks above for the first line of input, 
65.      * the dX calculation is the sum of 1 * leftData, 0 * middleData, and -1 * rightData. 
66.      * The dY calculation is the sum of 1 * leftData, 2 * middleData, and 1 * rightData. 
67.      * This is what is being calculated below, except we have removed the 
68.      * unnecessary calculations (multiplications by 1 or 0) and we are calculating 16 values at once. 
69.      * This pattern repeats for the other 2 rows of data. 
70.      */  
71.     short16 dx = rightData - leftData;  
72.     short16 dy = rightData + leftData + middleData * (short)2;  
73.     /* [Calculation] */  
74.   
75.     /* 
76.      * Second row of input. 
77.      * By adding the 'width * 1' to the offset we get the next row of data at the same column position. 
78.      * middleData is not loaded because it is not used in any of the calculations. 
79.      */  
80.     leftLoad = vload16(0, inputImage + (offset + width * 1 + 0));  
81.     rightLoad = vload16(0, inputImage + (offset + width * 1 + 2));  
82.   
83.     leftData = convert_short16(leftLoad);  
84.     rightData = convert_short16(rightLoad);  
85.   
86.     /* 
87.      * Calculate the results for the second row. 
88.      * The dX calculation is the sum of -2 * leftData, 0 * middleData, and -2 * rightData. 
89.      * There is no dY calculation to do: sum of 0 * leftData, 0 * middleData, and 0 * rightData. 
90.      */  
91.     dx += (rightData - leftData) * (short)2;  
92.   
93.     /* Third row of input. */  
94.     leftLoad = vload16(0, inputImage + (offset + width * 2 + 0));  
95.     middleLoad = vload16(0, inputImage + (offset + width * 2 + 1));  
96.     rightLoad = vload16(0, inputImage + (offset + width * 2 + 2));  
97.   
98.     leftData = convert_short16(leftLoad);  
99.     middleData = convert_short16(middleLoad);  
100.     rightData = convert_short16(rightLoad);  
101.   
102.     /* 
103.      * Calculate the results for the third row. 
104.      * The dX calculation is the sum of -1 * leftData, 0 * middleData, and -1 * rightData. 
105.      * The dY calculation is the sum of -1 * leftData, -2 * middleData, and -1 * rightData. 
106.      */  
107.     dx += rightData - leftData;  
108.     dy -= rightData + leftData + middleData * (short)2;  
109.   
110.     /* [Store] */  
111.     /* 
112.      * Store the results. 
113.      * The range of outputs from our Sobel calculations is [-1020, 1020]. 
114.      * In order to output this as an 8-bit signed char we must divide it by 8 (or shift right 3 times). 
115.      * This gives the range [-128, 128]. Depending on what type of output you require, 
116.      * (signed/unsigned, seperate/combined gradients) it is possible to do more of the calculations on the GPU using OpenCL. 
117.      * In this sample we're assuming that the application requires signed uncombined gradient outputs. 
118.      */  
119.     vstore16(convert_char16(dx >> 3), 0, outputImageDX + offset + width + 1);  
120.     vstore16(convert_char16(dy >> 3), 0, outputImageDY + offset + width + 1);  
121.     /* [Store] */  
122. }  

 

附录2：宿主机源码

[cpp]  view plain copy print ?

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.   
23. /** 
24.  * \brief Simple Sobel filter OpenCL sample. 
25.  * \details A sample which loads a bitmap and then passes it to the GPU. 
26.  *          An OpenCL kernel which does Sobel filtering is then run on the data. 
27.  *          The gradients of the image in x and y directions are returned by the GPU 
28.  *          and are combined on the CPU to form the filtered data. 
29.  *          The input image is loaded from assets/input.bmp. The output gradients in X and Y, 
30.  *          as well as the combined gradient image are stored in output-dX.bmp, output-dY.bmp 
31.  *          and output.bmp respectively. 
32.  * \return The exit code of the application, non-zero if a problem occurred. 
33.  */  
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. }