Style fix to OpenCL SVG Filters
[WebKit-https.git] / Source / WebCore / platform / graphics / gpu / opencl / FilterContextOpenCL.cpp
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
33 namespace WebCore {
34
35 FilterContextOpenCL* FilterContextOpenCL::m_context = 0;
36 int FilterContextOpenCL::m_alreadyInitialized = 0;
37
38 FilterContextOpenCL* 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
91 OpenCLHandle FilterContextOpenCL::createOpenCLImage(IntSize paintSize)
92 {
93     FilterContextOpenCL* context = FilterContextOpenCL::context();
94
95     cl_image_format clImageFormat;
96     clImageFormat.image_channel_order = CL_RGBA;
97     clImageFormat.image_channel_data_type = CL_UNORM_INT8;
98
99     OpenCLHandle image = clCreateImage2D(context->deviceContext(), CL_MEM_READ_WRITE, &clImageFormat,
100         paintSize.width(), paintSize.height(), 0, 0, 0);
101     return image;
102 }
103
104 static const char* transformColorSpaceKernelProgram =
105 PROGRAM(
106 const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
107
108 __kernel void transformColorSpace(__read_only image2d_t source, __write_only image2d_t destination, __constant float *clLookUpTable)
109 {
110     int2 sourceCoord = (int2) (get_global_id(0), get_global_id(1));
111     float4 pixel = read_imagef(source, sampler, sourceCoord);
112
113     pixel = (float4) (clLookUpTable[(int)(round(pixel.x * 255))], clLookUpTable[(int)(round(pixel.y * 255))],
114         clLookUpTable[(int) (round(pixel.z * 255))], pixel.w);
115
116     write_imagef(destination, sourceCoord, pixel);
117 }
118 ); // End of OpenCL kernels
119
120 void FilterContextOpenCL::openCLTransformColorSpace(OpenCLHandle& source, IntRect sourceSize, ColorSpace srcColorSpace, ColorSpace dstColorSpace)
121 {
122     DEFINE_STATIC_LOCAL(OpenCLHandle, deviceRgbLUT, ());
123     DEFINE_STATIC_LOCAL(OpenCLHandle, linearRgbLUT, ());
124
125     if (srcColorSpace == dstColorSpace)
126         return;
127
128     if ((srcColorSpace != ColorSpaceLinearRGB && srcColorSpace != ColorSpaceDeviceRGB)
129         || (dstColorSpace != ColorSpaceLinearRGB && dstColorSpace != ColorSpaceDeviceRGB))
130         return;
131
132     FilterContextOpenCL* context = FilterContextOpenCL::context();
133     ASSERT(context);
134
135     OpenCLHandle destination = context->createOpenCLImage(sourceSize.size());
136
137     if (!m_transformColorSpaceProgram) {
138         m_transformColorSpaceProgram = compileProgram(transformColorSpaceKernelProgram);
139         ASSERT(m_transformColorSpaceProgram);
140         m_transformColorSpaceKernel = kernelByName(m_transformColorSpaceProgram, "transformColorSpace");
141         ASSERT(m_transformColorSpaceKernel);
142     }
143
144     RunKernel kernel(context, m_transformColorSpaceKernel, sourceSize.width(), sourceSize.height());
145     kernel.addArgument(source);
146     kernel.addArgument(destination);
147
148     if (dstColorSpace == ColorSpaceLinearRGB) {
149         if (!linearRgbLUT) {
150             Vector<float> lookUpTable;
151             for (unsigned i = 0; i < 256; i++) {
152                 float color = i  / 255.0f;
153                 color = (color <= 0.04045f ? color / 12.92f : pow((color + 0.055f) / 1.055f, 2.4f));
154                 color = std::max(0.0f, color);
155                 color = std::min(1.0f, color);
156                 lookUpTable.append((round(color * 255)) / 255);
157             }
158             linearRgbLUT = kernel.addArgument(lookUpTable.data(), sizeof(float) * 256);
159         } else
160             kernel.addArgument(linearRgbLUT);
161     } else if (dstColorSpace == ColorSpaceDeviceRGB) {
162         if (!deviceRgbLUT) {
163             Vector<float> lookUpTable;
164             for (unsigned i = 0; i < 256; i++) {
165                 float color = i / 255.0f;
166                 color = (powf(color, 1.0f / 2.4f) * 1.055f) - 0.055f;
167                 color = std::max(0.0f, color);
168                 color = std::min(1.0f, color);
169                 lookUpTable.append((round(color * 255)) / 255);
170             }
171             deviceRgbLUT = kernel.addArgument(lookUpTable.data(), sizeof(float) * 256);
172         } else
173             kernel.addArgument(deviceRgbLUT);
174     }
175
176     kernel.run();
177     source.clear();
178     source = destination;
179 }
180
181 cl_program FilterContextOpenCL::compileProgram(const char* source)
182 {
183     cl_program program;
184     cl_int errorNumber;
185
186     FilterContextOpenCL* context = FilterContextOpenCL::context();
187     ASSERT(context);
188
189     program = clCreateProgramWithSource(context->m_deviceContext, 1, (const char**) &source, 0, 0);
190     errorNumber = clBuildProgram(program, 0, 0, 0, 0, 0);
191     if (errorNumber)
192         return 0;
193
194     return program;
195 }
196 } // namespace WebCore
197
198 #endif