1/*
2 * Copyright (C) 2012 University of Szeged
3 * Copyright (C) 2012 Tamas Czene <tczene@inf.u-szeged.hu>
4 * All rights reserved.
5 *
6 * Redistribution and use in source and binary forms, with or without
7 * modification, are permitted provided that the following conditions
8 * are met:
9 * 1. Redistributions of source code must retain the above copyright
10 *    notice, this list of conditions and the following disclaimer.
11 * 2. Redistributions in binary form must reproduce the above copyright
12 *    notice, this list of conditions and the following disclaimer in the
13 *    documentation and/or other materials provided with the distribution.
14 *
15 * THIS SOFTWARE IS PROVIDED BY UNIVERSITY OF SZEGED ``AS IS'' AND ANY
16 * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
17 * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
18 * PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL UNIVERSITY OF SZEGED OR
19 * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
20 * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
21 * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
22 * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
23 * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
24 * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
25 * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
26 */
27
28#include "config.h"
29
30#if ENABLE(OPENCL)
31#include "FilterContextOpenCL.h"
32
33namespace WebCore {
34
35FilterContextOpenCL* FilterContextOpenCL::m_context = 0;
36int FilterContextOpenCL::m_alreadyInitialized = 0;
37
38FilterContextOpenCL* FilterContextOpenCL::context()
39{
40    if (m_context)
41        return m_context;
42    if (m_alreadyInitialized)
43        return 0;
44
45    m_alreadyInitialized = true;
46    FilterContextOpenCL* localContext = new FilterContextOpenCL();
47
48    // Initializing the context.
49    cl_int errorNumber;
50    cl_device_id* devices;
51    cl_platform_id firstPlatformId;
52    size_t deviceBufferSize = 0;
53
54    errorNumber = clGetPlatformIDs(1, &firstPlatformId, 0);
55    cl_context_properties contextProperties[] = { CL_CONTEXT_PLATFORM, (cl_context_properties)firstPlatformId, 0};
56    localContext->m_deviceContext = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_GPU, 0, 0, &errorNumber);
57    if (errorNumber != CL_SUCCESS) {
58        localContext->m_deviceContext = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_CPU, 0, 0, &errorNumber);
59        if (errorNumber != CL_SUCCESS)
60            return 0;
61    }
62
63    errorNumber = clGetContextInfo(localContext->m_deviceContext, CL_CONTEXT_DEVICES, 0, 0, &deviceBufferSize);
64    if (errorNumber != CL_SUCCESS)
65        return 0;
66
67    if (!deviceBufferSize)
68        return 0;
69
70    devices = reinterpret_cast<cl_device_id*>(fastMalloc(deviceBufferSize));
71    errorNumber = clGetContextInfo(localContext->m_deviceContext, CL_CONTEXT_DEVICES, deviceBufferSize, devices, 0);
72    if (errorNumber != CL_SUCCESS)
73        return 0;
74
75    localContext->m_commandQueue = clCreateCommandQueue(localContext->m_deviceContext, devices[0], 0, 0);
76    if (!localContext->m_commandQueue)
77        return 0;
78
79    localContext->m_deviceId = devices[0];
80    fastFree(devices);
81
82    cl_bool imageSupport = CL_FALSE;
83    clGetDeviceInfo(localContext->m_deviceId, CL_DEVICE_IMAGE_SUPPORT, sizeof(cl_bool), &imageSupport, 0);
84    if (imageSupport != CL_TRUE)
85        return 0;
86
87    m_context = localContext;
88    return m_context;
89}
90
91void FilterContextOpenCL::freeResources()
92{
93    clFinish(m_commandQueue);
94
95    if (m_colorMatrixWasCompiled) {
96        freeResource(m_matrixOperation);
97        freeResource(m_saturateAndHueRotateOperation);
98        freeResource(m_luminanceOperation);
99        freeResource(m_colorMatrixProgram);
100    }
101    m_colorMatrixWasCompiled = false;
102
103    if (m_turbulenceWasCompiled) {
104        freeResource(m_turbulenceOperation);
105        freeResource(m_turbulenceProgram);
106    }
107    m_turbulenceWasCompiled = false;
108
109    if (m_transformColorSpaceWasCompiled) {
110        freeResource(m_transformColorSpaceKernel);
111        freeResource(m_transformColorSpaceProgram);
112    }
113    m_transformColorSpaceWasCompiled = false;
114}
115
116void FilterContextOpenCL::destroyContext()
117{
118    freeResources();
119
120    if (m_commandQueue)
121        clReleaseCommandQueue(m_commandQueue);
122    m_commandQueue = 0;
123
124    if (m_deviceContext)
125        clReleaseContext(m_deviceContext);
126    m_deviceContext = 0;
127
128    m_context = 0;
129}
130
131OpenCLHandle FilterContextOpenCL::createOpenCLImage(IntSize paintSize)
132{
133    FilterContextOpenCL* context = FilterContextOpenCL::context();
134
135    cl_image_format clImageFormat;
136    clImageFormat.image_channel_order = CL_RGBA;
137    clImageFormat.image_channel_data_type = CL_UNORM_INT8;
138
139#ifdef CL_API_SUFFIX__VERSION_1_2
140    cl_image_desc imageDescriptor = { CL_MEM_OBJECT_IMAGE2D, paintSize.width(), paintSize.height(), 0, 0, 0, 0, 0, 0, 0};
141    OpenCLHandle image = clCreateImage(context->deviceContext(), CL_MEM_READ_WRITE, &clImageFormat, &imageDescriptor, 0, 0);
142#else
143    OpenCLHandle image = clCreateImage2D(context->deviceContext(), CL_MEM_READ_WRITE, &clImageFormat,
144        paintSize.width(), paintSize.height(), 0, 0, 0);
145#endif
146    return image;
147}
148
149static const char* transformColorSpaceKernelProgram =
150PROGRAM(
151const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
152
153__kernel void transformColorSpace(__read_only image2d_t source, __write_only image2d_t destination, __constant float *clLookUpTable)
154{
155    int2 sourceCoord = (int2) (get_global_id(0), get_global_id(1));
156    float4 pixel = read_imagef(source, sampler, sourceCoord);
157
158    pixel = (float4) (clLookUpTable[(int)(round(pixel.x * 255))], clLookUpTable[(int)(round(pixel.y * 255))],
159        clLookUpTable[(int) (round(pixel.z * 255))], pixel.w);
160
161    write_imagef(destination, sourceCoord, pixel);
162}
163); // End of OpenCL kernels
164
165inline bool FilterContextOpenCL::compileTransformColorSpaceProgram()
166{
167    if (m_transformColorSpaceWasCompiled || inError())
168        return !inError();
169
170    m_transformColorSpaceWasCompiled = true;
171
172    if (isResourceAllocationFailed((m_transformColorSpaceProgram = compileProgram(transformColorSpaceKernelProgram))))
173        return false;
174    if (isResourceAllocationFailed((m_transformColorSpaceKernel = kernelByName(m_transformColorSpaceProgram, "transformColorSpace"))))
175        return false;
176    return true;
177}
178
179void FilterContextOpenCL::openCLTransformColorSpace(OpenCLHandle& source, IntRect sourceSize, ColorSpace srcColorSpace, ColorSpace dstColorSpace)
180{
181    DEFINE_STATIC_LOCAL(OpenCLHandle, deviceRgbLUT, ());
182    DEFINE_STATIC_LOCAL(OpenCLHandle, linearRgbLUT, ());
183
184    if (srcColorSpace == dstColorSpace || inError())
185        return;
186
187    if ((srcColorSpace != ColorSpaceLinearRGB && srcColorSpace != ColorSpaceDeviceRGB)
188        || (dstColorSpace != ColorSpaceLinearRGB && dstColorSpace != ColorSpaceDeviceRGB))
189        return;
190
191    if (!compileTransformColorSpaceProgram())
192        return;
193
194    OpenCLHandle destination = createOpenCLImage(sourceSize.size());
195
196    RunKernel kernel(this, m_transformColorSpaceKernel, sourceSize.width(), sourceSize.height());
197    kernel.addArgument(source);
198    kernel.addArgument(destination);
199
200    if (dstColorSpace == ColorSpaceLinearRGB) {
201        if (!linearRgbLUT) {
202            Vector<float> lookUpTable;
203            for (unsigned i = 0; i < 256; i++) {
204                float color = i  / 255.0f;
205                color = (color <= 0.04045f ? color / 12.92f : pow((color + 0.055f) / 1.055f, 2.4f));
206                color = std::max(0.0f, color);
207                color = std::min(1.0f, color);
208                lookUpTable.append((round(color * 255)) / 255);
209            }
210            linearRgbLUT = kernel.addArgument(lookUpTable.data(), sizeof(float) * 256);
211        } else
212            kernel.addArgument(linearRgbLUT);
213    } else if (dstColorSpace == ColorSpaceDeviceRGB) {
214        if (!deviceRgbLUT) {
215            Vector<float> lookUpTable;
216            for (unsigned i = 0; i < 256; i++) {
217                float color = i / 255.0f;
218                color = (powf(color, 1.0f / 2.4f) * 1.055f) - 0.055f;
219                color = std::max(0.0f, color);
220                color = std::min(1.0f, color);
221                lookUpTable.append((round(color * 255)) / 255);
222            }
223            deviceRgbLUT = kernel.addArgument(lookUpTable.data(), sizeof(float) * 256);
224        } else
225            kernel.addArgument(deviceRgbLUT);
226    }
227
228    kernel.run();
229    source.clear();
230    source = destination;
231}
232
233static const char* fillKernelProgram =
234PROGRAM_STR(
235__kernel void fill(__write_only image2d_t destination, float r, float g, float b, float a)
236{
237    float4 sourcePixel = (float4)(r, g, b, a);
238    write_imagef(destination, (int2)(get_global_id(0), get_global_id(1)), sourcePixel);
239}
240);
241
242inline bool FilterContextOpenCL::compileFill()
243{
244    if (m_fillWasCompiled || inError())
245        return !inError();
246
247    m_fillWasCompiled = true;
248
249    if (isResourceAllocationFailed((m_fillProgram = compileProgram(fillKernelProgram))))
250        return false;
251    if (isResourceAllocationFailed((m_fill = kernelByName(m_fillProgram, "fill"))))
252        return false;
253    return true;
254}
255
256void FilterContextOpenCL::fill(cl_mem image, IntSize imageSize, Color color)
257{
258    if (!m_context || inError())
259        return;
260
261    compileFill();
262
263    float r, g, b, a;
264
265    color.getRGBA(r, g, b, a);
266
267    RunKernel kernel(this, m_fill, imageSize.width(), imageSize.height());
268    kernel.addArgument(image);
269    kernel.addArgument(r);
270    kernel.addArgument(g);
271    kernel.addArgument(b);
272    kernel.addArgument(a);
273    kernel.run();
274}
275
276cl_program FilterContextOpenCL::compileProgram(const char* source)
277{
278    cl_program program;
279    cl_int errorNumber = 0;
280
281    program = clCreateProgramWithSource(m_deviceContext, 1, (const char**) &source, 0, &errorNumber);
282    if (isFailed(errorNumber))
283        return 0;
284
285    if (isFailed(clBuildProgram(program, 0, 0, 0, 0, 0)))
286        return 0;
287
288    return program;
289}
290
291void FilterContextOpenCL::freeResource(cl_kernel& handle)
292{
293    if (handle) {
294        clReleaseKernel(handle);
295        handle = 0;
296    }
297}
298
299void FilterContextOpenCL::freeResource(cl_program& handle)
300{
301    if (handle) {
302        clReleaseProgram(handle);
303        handle = 0;
304    }
305}
306} // namespace WebCore
307
308#endif
309