Style fix to OpenCL SVG Filters
[WebKit-https.git] / Source / WebCore / platform / graphics / gpu / opencl / OpenCLFEColorMatrix.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(FILTERS) && ENABLE(OPENCL)
31 #include "FEColorMatrix.h"
32
33 #include "FilterContextOpenCL.h"
34
35 namespace WebCore {
36
37 #define COLOR_MATRIX_KERNEL(...) \
38         int2 sourceCoord = (int2) (get_global_id(0) + x, get_global_id(1) + y); \
39         float4 sourcePixel = read_imagef(source, sampler, sourceCoord); \
40         float4 destinationPixel = (float4) (__VA_ARGS__); \
41         write_imagef(destination, (int2) (get_global_id(0), get_global_id(1)), destinationPixel);
42
43 static const char* colorMatrixKernelProgram =
44 PROGRAM(
45 const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
46
47 __kernel void matrix(__write_only image2d_t destination, __read_only image2d_t source, float x, float y, __constant float *values)
48 {
49     COLOR_MATRIX_KERNEL(values[0] * sourcePixel.x + values[1] * sourcePixel.y + values[2] * sourcePixel.z + values[3] * sourcePixel.w + values[4],
50         values[5] * sourcePixel.x + values[6] * sourcePixel.y + values[7] * sourcePixel.z + values[8] * sourcePixel.w + values[9],
51         values[10] * sourcePixel.x + values[11] * sourcePixel.y + values[12] * sourcePixel.z + values[13] * sourcePixel.w + values[14],
52         values[15] * sourcePixel.x + values[16] * sourcePixel.y + values[17] * sourcePixel.z + values[18] * sourcePixel.w + values[19])
53 }
54
55 __kernel void saturateAndHueRotate(__write_only image2d_t destination, __read_only image2d_t source, float x, float y, __constant float *components)
56 {
57     COLOR_MATRIX_KERNEL(sourcePixel.x * components[0] + sourcePixel.y * components[1] + sourcePixel.z * components[2],
58         sourcePixel.x * components[3] + sourcePixel.y * components[4] + sourcePixel.z * components[5],
59         sourcePixel.x * components[6] + sourcePixel.y * components[7] + sourcePixel.z * components[8],
60         sourcePixel.w)
61 }
62
63 __kernel void luminance(__write_only image2d_t destination, __read_only image2d_t source, float x, float y)
64 {
65     COLOR_MATRIX_KERNEL(0, 0, 0, 0.2125 * sourcePixel.x + 0.7154 * sourcePixel.y + 0.0721 * sourcePixel.z)
66 }
67 ); // End of OpenCL kernels
68
69 inline bool FilterContextOpenCL::compileFEColorMatrix()
70 {
71     if (m_colorMatrixCompileStatus != openclNotCompiledYet)
72         return m_colorMatrixCompileStatus == openclCompileSuccessful;
73
74     m_colorMatrixCompileStatus = openclCompileFailed;
75     m_colorMatrixProgram = compileProgram(colorMatrixKernelProgram);
76     if (!m_colorMatrixProgram)
77         return false;
78
79     m_matrixOperation = kernelByName(m_colorMatrixProgram, "matrix");
80     if (!m_matrixOperation)
81         return false;
82     m_saturateAndHueRotateOperation = kernelByName(m_colorMatrixProgram, "saturateAndHueRotate");
83     if (!m_saturateAndHueRotateOperation)
84         return false;
85     m_luminanceOperation = kernelByName(m_colorMatrixProgram, "luminance");
86     if (!m_luminanceOperation)
87         return false;
88
89     m_colorMatrixCompileStatus = openclCompileSuccessful;
90     return openclCompileSuccessful;
91 }
92
93 inline void FilterContextOpenCL::applyFEColorMatrix(OpenCLHandle destination, IntSize destinationSize, OpenCLHandle source, IntPoint relativeSourceLocation, float* values, int type)
94 {
95     cl_kernel colorMatrix;
96     OpenCLHandle clValues;
97
98     switch (type) {
99     case FECOLORMATRIX_TYPE_MATRIX:
100         colorMatrix = m_matrixOperation;
101         break;
102     case FECOLORMATRIX_TYPE_SATURATE:
103         colorMatrix = m_saturateAndHueRotateOperation;
104         break;
105     case FECOLORMATRIX_TYPE_HUEROTATE:
106         colorMatrix = m_saturateAndHueRotateOperation;
107         break;
108     case FECOLORMATRIX_TYPE_LUMINANCETOALPHA:
109         colorMatrix = m_luminanceOperation;
110         break;
111     default:
112         ASSERT_NOT_REACHED();
113         return;
114     }
115
116     RunKernel kernel(this, colorMatrix, destinationSize.width(), destinationSize.height());
117     kernel.addArgument(destination);
118     kernel.addArgument(source);
119     kernel.addArgument(relativeSourceLocation.x());
120     kernel.addArgument(relativeSourceLocation.y());
121     if (type == FECOLORMATRIX_TYPE_MATRIX)
122         clValues = kernel.addArgument(values, sizeof(float) * 20);
123     else if (type == FECOLORMATRIX_TYPE_SATURATE || type == FECOLORMATRIX_TYPE_HUEROTATE)
124         clValues = kernel.addArgument(values, sizeof(float) * 9);
125     kernel.run();
126
127     clValues.clear();
128 }
129
130 bool FEColorMatrix::platformApplyOpenCL()
131 {
132     FilterContextOpenCL* context = FilterContextOpenCL::context();
133     if (!context || !context->compileFEColorMatrix())
134         return false;
135
136     FilterEffect* in = inputEffect(0);
137     OpenCLHandle source = in->openCLImage();
138     OpenCLHandle destination = createOpenCLImageResult();
139
140     IntPoint relativeSourceLocation(
141         absolutePaintRect().x() - in->absolutePaintRect().location().x(),
142         absolutePaintRect().y() - in->absolutePaintRect().location().y());
143
144     float components[9];
145     if (FECOLORMATRIX_TYPE_SATURATE == m_type)
146         calculateSaturateComponents(components, m_values[0]);
147     else if (FECOLORMATRIX_TYPE_HUEROTATE == m_type)
148         calculateHueRotateComponents(components, m_values[0]);
149
150     context->applyFEColorMatrix(destination, absolutePaintRect().size(), source, relativeSourceLocation, (FECOLORMATRIX_TYPE_MATRIX == m_type) ? m_values.data() : components, m_type);
151
152     return true;
153 }
154
155 } // namespace WebCore
156
157 #endif // ENABLE(FILTERS) && ENABLE(OPENCL)