7ebc6074d1015921a2a0ba65e3a1f3ab0029c71f
[WebKit-https.git] / Source / WebCore / platform / graphics / gpu / opencl / OpenCLFETurbulence.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 #include "config.h"
28
29 #if ENABLE(FILTERS) && ENABLE(OPENCL)
30 #include "FETurbulence.h"
31
32 #include "FETurbulence.cpp"
33 #include "FilterContextOpenCL.h"
34 #include "SVGFilter.h"
35
36 namespace WebCore {
37
38 static const char* turbulenceKernelProgram =
39 PROGRAM_STR(
40 __constant int s_perlinNoise = 4096;
41 __constant int s_blockSize = 256;
42 __constant int s_blockMask = 255;
43
44 typedef struct {
45     int noisePositionIntegerValue;
46     float noisePositionFractionValue;
47 } Noise;
48
49 typedef struct {
50     int width;
51     int wrapX;
52     int height;
53     int wrapY;
54 } StitchData;
55
56 float linearInterpolation(float t, float a, float b)
57 {
58     return mad(b - a, t, a);
59 }
60
61 float noise2D(__constant float *component, __constant int *latticeSelector, StitchData stitchData, float noiseVectorX, float noiseVectorY, int stitchTiles)
62 {
63     Noise noiseX;
64     noiseX.noisePositionIntegerValue = (int)(noiseVectorX + s_perlinNoise);
65     noiseX.noisePositionFractionValue = (noiseVectorX + s_perlinNoise) - noiseX.noisePositionIntegerValue;
66     Noise noiseY;
67     noiseY.noisePositionIntegerValue = (int)(noiseVectorY + s_perlinNoise);
68     noiseY.noisePositionFractionValue = (noiseVectorY + s_perlinNoise) - noiseY.noisePositionIntegerValue;
69
70     // If stitching, adjust lattice points accordingly.
71     if (stitchTiles) {
72         if (noiseX.noisePositionIntegerValue >= stitchData.wrapX)
73             noiseX.noisePositionIntegerValue -= stitchData.width;
74         if (noiseX.noisePositionIntegerValue >= stitchData.wrapX - 1)
75             noiseX.noisePositionIntegerValue -= stitchData.width - 1;
76         if (noiseY.noisePositionIntegerValue >= stitchData.wrapY)
77             noiseY.noisePositionIntegerValue -= stitchData.height;
78         if (noiseY.noisePositionIntegerValue >= stitchData.wrapY - 1)
79             noiseY.noisePositionIntegerValue -= stitchData.height - 1;
80     }
81
82     noiseX.noisePositionIntegerValue &= s_blockMask;
83     noiseY.noisePositionIntegerValue &= s_blockMask;
84     int latticeIndex = latticeSelector[noiseX.noisePositionIntegerValue];
85     int nextLatticeIndex = latticeSelector[(noiseX.noisePositionIntegerValue + 1) & s_blockMask];
86
87     float sx = noiseX.noisePositionFractionValue * noiseX.noisePositionFractionValue * (3 - 2 * noiseX.noisePositionFractionValue);
88     float sy = noiseY.noisePositionFractionValue * noiseY.noisePositionFractionValue * (3 - 2 * noiseY.noisePositionFractionValue);
89
90     // This is taken 1:1 from SVG spec: http://www.w3.org/TR/SVG11/filters.html#feTurbulenceElement.
91     int temp = latticeSelector[latticeIndex + noiseY.noisePositionIntegerValue];
92     float u = noiseX.noisePositionFractionValue * component[temp * 2] + noiseY.noisePositionFractionValue * component[temp * 2 + 1];
93     temp = latticeSelector[nextLatticeIndex + noiseY.noisePositionIntegerValue];
94     float v = (noiseX.noisePositionFractionValue - 1) * component[temp * 2] + noiseY.noisePositionFractionValue * component[temp * 2 + 1];
95     float a = linearInterpolation(sx, u, v);
96     temp = latticeSelector[latticeIndex + noiseY.noisePositionIntegerValue + 1];
97     u = noiseX.noisePositionFractionValue * component[temp * 2] + (noiseY.noisePositionFractionValue - 1) * component[temp * 2 + 1];
98     temp = latticeSelector[nextLatticeIndex + noiseY.noisePositionIntegerValue + 1];
99     v = (noiseX.noisePositionFractionValue - 1) * component[temp * 2] + (noiseY.noisePositionFractionValue - 1) * component[temp * 2 + 1];
100     float b = linearInterpolation(sx, u, v);
101     return linearInterpolation(sy, a, b);
102 }
103
104 __kernel void Turbulence(__write_only image2d_t destination, __constant float *transform, __constant float *redComponent,
105     __constant float *greenComponent, __constant float *blueComponent, __constant float *alphaComponent,
106     __constant int *latticeSelector, __private int offsetX, __private int offsetY, __private int tileWidth,
107     __private int tileHeight, __private float baseFrequencyX, __private float baseFrequencyY, __private int stitchTiles,
108     __private int numOctaves, __private int type, __private int filter_height)
109 {
110     StitchData stitchData = { 0, 0, 0, 0 };
111     // Adjust the base frequencies if necessary for stitching.
112     if (stitchTiles) {
113         // When stitching tiled turbulence, the frequencies must be adjusted
114         // so that the tile borders will be continuous.
115         if (baseFrequencyX) {
116             float lowFrequency = floor(tileWidth * baseFrequencyX) / tileWidth;
117             float highFrequency = ceil(tileWidth * baseFrequencyX) / tileWidth;
118             // BaseFrequency should be non-negative according to the standard.
119             baseFrequencyX = (baseFrequencyX / lowFrequency < highFrequency / baseFrequencyX) ? lowFrequency : highFrequency;
120         }
121         if (baseFrequencyY) {
122             float lowFrequency = floor(tileHeight * baseFrequencyY) / tileHeight;
123             float highFrequency = ceil(tileHeight * baseFrequencyY) / tileHeight;
124             baseFrequencyY = (baseFrequencyY / lowFrequency < highFrequency / baseFrequencyY) ? lowFrequency : highFrequency;
125         }
126         // Set up TurbulenceInitial stitch values.
127         stitchData.width = round(tileWidth * baseFrequencyX);
128         stitchData.wrapX = s_perlinNoise + stitchData.width;
129         stitchData.height = round(tileHeight * baseFrequencyY);
130         stitchData.wrapY = s_perlinNoise + stitchData.height;
131     }
132     float4 turbulenceFunctionResult = (float4)(0, 0, 0, 0);
133     float x = (get_global_id(0) + offsetX) * baseFrequencyX;
134     float y = (get_global_id(1) + offsetY) * baseFrequencyY;
135
136     float noiseVectorX = transform[0] * x + transform[2] * y + transform[4];
137     float noiseVectorY = transform[1] * x + transform[3] * y + transform[5];
138
139     float ratio = 1;
140     for (int octave = 0; octave < numOctaves; ++octave) {
141         float4 noise2DResult = (float4)( noise2D(redComponent, latticeSelector, stitchData, noiseVectorX, noiseVectorY, stitchTiles) / ratio,
142     noise2D(greenComponent, latticeSelector, stitchData, noiseVectorX, noiseVectorY, stitchTiles) / ratio,
143     noise2D(blueComponent, latticeSelector, stitchData, noiseVectorX, noiseVectorY, stitchTiles) / ratio,
144     noise2D(alphaComponent, latticeSelector, stitchData, noiseVectorX, noiseVectorY, stitchTiles) / ratio);
145
146         turbulenceFunctionResult += (type == 1) ? noise2DResult : fabs(noise2DResult);
147
148         noiseVectorX *= 2;
149         noiseVectorY *= 2;
150         ratio *= 2;
151         if (stitchTiles) {
152             // Update stitch values. Subtracting s_perlinNoiseoise before the multiplication and
153             // adding it afterward simplifies to subtracting it once.
154             stitchData.width *= 2;
155             stitchData.wrapX = 2 * stitchData.wrapX - s_perlinNoise;
156             stitchData.height *= 2;
157             stitchData.wrapY = 2 * stitchData.wrapY - s_perlinNoise;
158         }
159     }
160
161     if (type == 1)
162         turbulenceFunctionResult = mad(0.5f, turbulenceFunctionResult, 0.5f);
163     // Clamp result.
164     turbulenceFunctionResult = clamp(turbulenceFunctionResult, 0.0f, 1.0f);
165
166     write_imagef(destination, (int2)(get_global_id(0), get_global_id(1)), turbulenceFunctionResult);
167 }
168 );
169
170 inline bool FilterContextOpenCL::compileFETurbulence()
171 {
172     if (m_turbulenceCompileStatus != openclNotCompiledYet)
173         return m_turbulenceCompileStatus == openclCompileSuccessful;
174
175     m_turbulenceCompileStatus = openclCompileFailed;
176     m_turbulenceProgram = compileProgram(turbulenceKernelProgram);
177     if (!m_turbulenceProgram)
178         return false;
179     m_turbulenceOperation = kernelByName(m_turbulenceProgram, "Turbulence");
180     if (!m_turbulenceOperation)
181         return false;
182
183     m_turbulenceCompileStatus = openclCompileSuccessful;
184     return openclCompileFailed;
185 }
186
187 inline void FilterContextOpenCL::applyFETurbulence(OpenCLHandle destination,
188     IntSize destinationSize, int blockSize,
189     void* transform, void* redComponent, void* greenComponent,
190     void* blueComponent, void* alphaComponent,
191     void* latticeSelector, int offsetX, int offsetY, int tileWidth, int tileHeight,
192     float baseFrequencyX, float baseFrequencyY, bool stitchTiles, int numOctaves, int type)
193 {
194     RunKernel kernel(this, m_turbulenceOperation, destinationSize.width(), destinationSize.height());
195
196     kernel.addArgument(destination);
197     OpenCLHandle transformHandle(kernel.addArgument(transform, sizeof(float) * 6));
198     OpenCLHandle redComponentHandle(kernel.addArgument(redComponent, sizeof(float) * (2 * blockSize + 2) * 2));
199     OpenCLHandle greenComponentHandle(kernel.addArgument(greenComponent, sizeof(float) * (2 * blockSize + 2) * 2));
200     OpenCLHandle blueComponentHandle(kernel.addArgument(blueComponent, sizeof(float) * (2 * blockSize + 2) * 2));
201     OpenCLHandle alphaComponentHandle(kernel.addArgument(alphaComponent, sizeof(float) * (2 * blockSize + 2) * 2));
202     OpenCLHandle latticeSelectorHandle(kernel.addArgument(latticeSelector, sizeof(int) * (2 * blockSize + 2)));
203     kernel.addArgument(offsetX);
204     kernel.addArgument(offsetY);
205     kernel.addArgument(tileWidth);
206     kernel.addArgument(tileHeight);
207     kernel.addArgument(baseFrequencyX);
208     kernel.addArgument(baseFrequencyY);
209     kernel.addArgument(stitchTiles);
210     kernel.addArgument(numOctaves);
211     kernel.addArgument(type);
212     kernel.addArgument(destinationSize.height());
213
214     kernel.run();
215
216     transformHandle.clear();
217     redComponentHandle.clear();
218     greenComponentHandle.clear();
219     blueComponentHandle.clear();
220     alphaComponentHandle.clear();
221     latticeSelectorHandle.clear();
222 }
223
224 bool FETurbulence::platformApplyOpenCL()
225 {
226     FilterContextOpenCL* context = FilterContextOpenCL::context();
227     if (!context || !context->compileFETurbulence())
228         return false;
229
230     OpenCLHandle destination = createOpenCLImageResult();
231
232     PaintingData paintingData(m_seed, roundedIntSize(filterPrimitiveSubregion().size()));
233     initPaint(paintingData);
234
235     AffineTransform invertedTransform = reinterpret_cast<SVGFilter*>(filter())->absoluteTransform().inverse();
236     float transformComponents[6] = { invertedTransform.a(), invertedTransform.b(), invertedTransform.c(), invertedTransform.d(), invertedTransform.e(), invertedTransform.f() };
237
238     context->applyFETurbulence(destination, absolutePaintRect().size(), s_blockSize, transformComponents, paintingData.gradient,
239         paintingData.gradient + 1, paintingData.gradient + 2, paintingData.gradient + 3, paintingData.latticeSelector,
240         absolutePaintRect().x(), absolutePaintRect().y(), paintingData.filterSize.width(), paintingData.filterSize.height(),
241         m_baseFrequencyX, m_baseFrequencyY, m_stitchTiles, m_numOctaves, m_type);
242
243     return true;
244 }
245
246 } // namespace WebCore
247
248 #endif // ENABLE(FILTERS) && ENABLE(OPENCL)