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
36namespace WebCore {
37
38static const char* turbulenceKernelProgram =
39PROGRAM(
40__constant int s_perlinNoise = 4096;
41__constant int s_blockSize = 256;
42__constant int s_blockMask = 255;
43
44typedef struct {
45    int noisePositionIntegerValue;
46    float noisePositionFractionValue;
47} Noise;
48
49typedef struct {
50    int width;
51    int wrapX;
52    int height;
53    int wrapY;
54} StitchData;
55
56float linearInterpolation(float t, float a, float b)
57{
58    return mad(b - a, t, a);
59}
60
61float 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); // End of OpenCL kernels
169
170inline bool FilterContextOpenCL::compileFETurbulence()
171{
172    if (m_turbulenceWasCompiled || inError())
173        return !inError();
174
175    m_turbulenceWasCompiled = true;
176
177    if (isResourceAllocationFailed((m_turbulenceProgram = compileProgram(turbulenceKernelProgram))))
178        return false;
179    if (isResourceAllocationFailed((m_turbulenceOperation = kernelByName(m_turbulenceProgram, "Turbulence"))))
180        return false;
181    return true;
182}
183
184inline void FilterContextOpenCL::applyFETurbulence(OpenCLHandle destination,
185    IntSize destinationSize, int totalBlockSize,
186    void* transform, void* redComponent, void* greenComponent,
187    void* blueComponent, void* alphaComponent,
188    int* latticeSelector, int offsetX, int offsetY, int tileWidth, int tileHeight,
189    float baseFrequencyX, float baseFrequencyY, bool stitchTiles, int numOctaves, int type)
190{
191    RunKernel kernel(this, m_turbulenceOperation, destinationSize.width(), destinationSize.height());
192
193    kernel.addArgument(destination);
194    OpenCLHandle transformHandle(kernel.addArgument(transform, sizeof(float) * 6));
195    OpenCLHandle redComponentHandle(kernel.addArgument(redComponent, sizeof(float) * totalBlockSize * 2));
196    OpenCLHandle greenComponentHandle(kernel.addArgument(greenComponent, sizeof(float) * totalBlockSize * 2));
197    OpenCLHandle blueComponentHandle(kernel.addArgument(blueComponent, sizeof(float) * totalBlockSize * 2));
198    OpenCLHandle alphaComponentHandle(kernel.addArgument(alphaComponent, sizeof(float) * totalBlockSize * 2));
199    OpenCLHandle latticeSelectorHandle(kernel.addArgument(latticeSelector, sizeof(int) * totalBlockSize));
200    kernel.addArgument(offsetX);
201    kernel.addArgument(offsetY);
202    kernel.addArgument(tileWidth);
203    kernel.addArgument(tileHeight);
204    kernel.addArgument(baseFrequencyX);
205    kernel.addArgument(baseFrequencyY);
206    kernel.addArgument(stitchTiles);
207    kernel.addArgument(numOctaves);
208    kernel.addArgument(type);
209    kernel.addArgument(destinationSize.height());
210
211    kernel.run();
212
213    transformHandle.clear();
214    redComponentHandle.clear();
215    greenComponentHandle.clear();
216    blueComponentHandle.clear();
217    alphaComponentHandle.clear();
218    latticeSelectorHandle.clear();
219}
220
221bool FETurbulence::platformApplyOpenCL()
222{
223    FilterContextOpenCL* context = FilterContextOpenCL::context();
224    if (!context)
225        return false;
226
227    if (!context->compileFETurbulence())
228        return true;
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(), 2 * s_blockSize + 2, 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)
249