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