Diff
Modified: trunk/Source/WebCore/ChangeLog (137590 => 137591)
--- trunk/Source/WebCore/ChangeLog 2012-12-13 15:04:36 UTC (rev 137590)
+++ trunk/Source/WebCore/ChangeLog 2012-12-13 15:06:39 UTC (rev 137591)
@@ -1,3 +1,29 @@
+2012-12-13 Tamas Czene <[email protected]>
+
+ OpenCL version of FEColorMatrix.
+ https://bugs.webkit.org/show_bug.cgi?id=103398
+
+ Reviewed by Zoltan Herczeg.
+
+ ~3x speed-up (depending on hardware configuration)
+
+ * Target.pri: Added a new file to the build system.
+ * platform/graphics/filters/FEColorMatrix.cpp: calculateSaturateComponents and calculateHueRotateComponents methodes are moved to the header, since they are used by OpenCL as well
+ (WebCore::effectType):
+ * platform/graphics/filters/FEColorMatrix.h:
+ (FEColorMatrix):
+ (WebCore::FEColorMatrix::calculateSaturateComponents):
+ (WebCore):
+ (WebCore::FEColorMatrix::calculateHueRotateComponents):
+ * platform/graphics/gpu/opencl/FilterContextOpenCL.h: Required fields for the ColorMatrix kernels were added
+ (WebCore::FilterContextOpenCL::FilterContextOpenCL):
+ (FilterContextOpenCL):
+ * platform/graphics/gpu/opencl/OpenCLFEColorMatrix.cpp: the core implementation goes here
+ (WebCore):
+ (WebCore::FilterContextOpenCL::compileFEColorMatrix):
+ (WebCore::FilterContextOpenCL::applyFEColorMatrix):
+ (WebCore::FEColorMatrix::platformApplyOpenCL):
+
2012-12-13 Ilya Tikhonovsky <[email protected]>
Web Inspector: Native Memory Instrumentation: do not validate pointers to objects in RenderArena agains tcmalloc data.
Modified: trunk/Source/WebCore/Target.pri (137590 => 137591)
--- trunk/Source/WebCore/Target.pri 2012-12-13 15:04:36 UTC (rev 137590)
+++ trunk/Source/WebCore/Target.pri 2012-12-13 15:06:39 UTC (rev 137591)
@@ -4107,6 +4107,7 @@
platform/graphics/gpu/opencl/FilterContextOpenCL.h
SOURCES += \
platform/graphics/gpu/opencl/FilterContextOpenCL.cpp \
+ platform/graphics/gpu/opencl/OpenCLFEColorMatrix.cpp \
platform/graphics/gpu/opencl/OpenCLFESourceAlpha.cpp \
platform/graphics/gpu/opencl/OpenCLFESourceGraphic.cpp \
platform/graphics/gpu/opencl/OpenCLFETurbulence.cpp
Modified: trunk/Source/WebCore/platform/graphics/filters/FEColorMatrix.cpp (137590 => 137591)
--- trunk/Source/WebCore/platform/graphics/filters/FEColorMatrix.cpp 2012-12-13 15:04:36 UTC (rev 137590)
+++ trunk/Source/WebCore/platform/graphics/filters/FEColorMatrix.cpp 2012-12-13 15:06:39 UTC (rev 137591)
@@ -105,34 +105,6 @@
blue = 0;
}
-inline void calculateSaturateComponents(float* components, float value)
-{
- components[0] = (0.213 + 0.787 * value);
- components[1] = (0.715 - 0.715 * value);
- components[2] = (0.072 - 0.072 * value);
- components[3] = (0.213 - 0.213 * value);
- components[4] = (0.715 + 0.285 * value);
- components[5] = (0.072 - 0.072 * value);
- components[6] = (0.213 - 0.213 * value);
- components[7] = (0.715 - 0.715 * value);
- components[8] = (0.072 + 0.928 * value);
-}
-
-inline void calculateHueRotateComponents(float* components, float value)
-{
- float cosHue = cos(value * piFloat / 180);
- float sinHue = sin(value * piFloat / 180);
- components[0] = 0.213 + cosHue * 0.787 - sinHue * 0.213;
- components[1] = 0.715 - cosHue * 0.715 - sinHue * 0.715;
- components[2] = 0.072 - cosHue * 0.072 + sinHue * 0.928;
- components[3] = 0.213 - cosHue * 0.213 + sinHue * 0.143;
- components[4] = 0.715 + cosHue * 0.285 + sinHue * 0.140;
- components[5] = 0.072 - cosHue * 0.072 - sinHue * 0.283;
- components[6] = 0.213 - cosHue * 0.213 - sinHue * 0.787;
- components[7] = 0.715 - cosHue * 0.715 + sinHue * 0.715;
- components[8] = 0.072 + cosHue * 0.928 + sinHue * 0.072;
-}
-
template<ColorMatrixType filterType>
void effectType(Uint8ClampedArray* pixelArray, const Vector<float>& values)
{
@@ -140,9 +112,9 @@
float components[9];
if (filterType == FECOLORMATRIX_TYPE_SATURATE)
- calculateSaturateComponents(components, values[0]);
+ FEColorMatrix::calculateSaturateComponents(components, values[0]);
else if (filterType == FECOLORMATRIX_TYPE_HUEROTATE)
- calculateHueRotateComponents(components, values[0]);
+ FEColorMatrix::calculateHueRotateComponents(components, values[0]);
for (unsigned pixelByteOffset = 0; pixelByteOffset < pixelArrayLength; pixelByteOffset += 4) {
float red = pixelArray->item(pixelByteOffset);
Modified: trunk/Source/WebCore/platform/graphics/filters/FEColorMatrix.h (137590 => 137591)
--- trunk/Source/WebCore/platform/graphics/filters/FEColorMatrix.h 2012-12-13 15:04:36 UTC (rev 137590)
+++ trunk/Source/WebCore/platform/graphics/filters/FEColorMatrix.h 2012-12-13 15:06:39 UTC (rev 137591)
@@ -53,10 +53,16 @@
virtual bool platformApplySkia();
virtual SkImageFilter* createImageFilter(SkiaImageFilterBuilder*);
#endif
+#if ENABLE(OPENCL)
+ virtual bool platformApplyOpenCL();
+#endif
virtual void dump();
virtual TextStream& externalRepresentation(TextStream&, int indention) const;
+ static inline void calculateSaturateComponents(float* components, float value);
+ static inline void calculateHueRotateComponents(float* components, float value);
+
private:
FEColorMatrix(Filter*, ColorMatrixType, const Vector<float>&);
@@ -64,6 +70,35 @@
Vector<float> m_values;
};
+inline void FEColorMatrix::calculateSaturateComponents(float* components, float value)
+{
+ components[0] = (0.213 + 0.787 * value);
+ components[1] = (0.715 - 0.715 * value);
+ components[2] = (0.072 - 0.072 * value);
+ components[3] = (0.213 - 0.213 * value);
+ components[4] = (0.715 + 0.285 * value);
+ components[5] = (0.072 - 0.072 * value);
+ components[6] = (0.213 - 0.213 * value);
+ components[7] = (0.715 - 0.715 * value);
+ components[8] = (0.072 + 0.928 * value);
+}
+
+inline void FEColorMatrix::calculateHueRotateComponents(float* components, float value)
+{
+ float cosHue = cos(value * piFloat / 180);
+ float sinHue = sin(value * piFloat / 180);
+ components[0] = 0.213 + cosHue * 0.787 - sinHue * 0.213;
+ components[1] = 0.715 - cosHue * 0.715 - sinHue * 0.715;
+ components[2] = 0.072 - cosHue * 0.072 + sinHue * 0.928;
+ components[3] = 0.213 - cosHue * 0.213 + sinHue * 0.143;
+ components[4] = 0.715 + cosHue * 0.285 + sinHue * 0.140;
+ components[5] = 0.072 - cosHue * 0.072 - sinHue * 0.283;
+ components[6] = 0.213 - cosHue * 0.213 - sinHue * 0.787;
+ components[7] = 0.715 - cosHue * 0.715 + sinHue * 0.715;
+ components[8] = 0.072 + cosHue * 0.928 + sinHue * 0.072;
+}
+
+
} // namespace WebCore
#endif // ENABLE(FILTERS)
Modified: trunk/Source/WebCore/platform/graphics/gpu/opencl/FilterContextOpenCL.h (137590 => 137591)
--- trunk/Source/WebCore/platform/graphics/gpu/opencl/FilterContextOpenCL.h 2012-12-13 15:04:36 UTC (rev 137590)
+++ trunk/Source/WebCore/platform/graphics/gpu/opencl/FilterContextOpenCL.h 2012-12-13 15:06:39 UTC (rev 137591)
@@ -49,6 +49,11 @@
, m_commandQueue(0)
, m_transformColorSpaceProgram(0)
, m_transformColorSpaceKernel(0)
+ , m_colorMatrixCompileStatus(openclNotCompiledYet)
+ , m_colorMatrixProgram(0)
+ , m_matrixOperation(0)
+ , m_saturateAndHueRotateOperation(0)
+ , m_luminanceOperation(0)
, m_turbulenceCompileStatus(openclNotCompiledYet)
, m_turbulenceProgram(0)
, m_turbulenceOperation(0)
@@ -65,8 +70,10 @@
OpenCLHandle createOpenCLImage(IntSize);
void openCLTransformColorSpace(OpenCLHandle&, IntRect, ColorSpace, ColorSpace);
+ inline bool compileFEColorMatrix();
inline bool compileFETurbulence();
+ inline void applyFEColorMatrix(OpenCLHandle, IntSize, OpenCLHandle, IntPoint, void*, int);
inline void applyFETurbulence(OpenCLHandle, IntSize, int, void*, void*, void*, void*, void*,
void*, int, int, int, int, float, float, bool, int, int);
@@ -142,6 +149,12 @@
cl_program m_transformColorSpaceProgram;
cl_kernel m_transformColorSpaceKernel;
+ OpenCLCompileStatus m_colorMatrixCompileStatus;
+ cl_program m_colorMatrixProgram;
+ cl_kernel m_matrixOperation;
+ cl_kernel m_saturateAndHueRotateOperation;
+ cl_kernel m_luminanceOperation;
+
OpenCLCompileStatus m_turbulenceCompileStatus;
cl_program m_turbulenceProgram;
cl_kernel m_turbulenceOperation;
Added: trunk/Source/WebCore/platform/graphics/gpu/opencl/OpenCLFEColorMatrix.cpp (0 => 137591)
--- trunk/Source/WebCore/platform/graphics/gpu/opencl/OpenCLFEColorMatrix.cpp (rev 0)
+++ trunk/Source/WebCore/platform/graphics/gpu/opencl/OpenCLFEColorMatrix.cpp 2012-12-13 15:06:39 UTC (rev 137591)
@@ -0,0 +1,157 @@
+/*
+ * Copyright (C) 2012 University of Szeged
+ * Copyright (C) 2012 Tamas Czene <[email protected]>
+ * All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions
+ * are met:
+ * 1. Redistributions of source code must retain the above copyright
+ * notice, this list of conditions and the following disclaimer.
+ * 2. Redistributions in binary form must reproduce the above copyright
+ * notice, this list of conditions and the following disclaimer in the
+ * documentation and/or other materials provided with the distribution.
+ *
+ * THIS SOFTWARE IS PROVIDED BY UNIVERSITY OF SZEGED ``AS IS'' AND ANY
+ * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL UNIVERSITY OF SZEGED OR
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
+ * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+
+#include "config.h"
+
+#if ENABLE(FILTERS) && ENABLE(OPENCL)
+#include "FEColorMatrix.h"
+
+#include "FilterContextOpenCL.h"
+
+namespace WebCore {
+
+#define COLOR_MATRIX_KERNEL(...) \
+ int2 sourceCoord = (int2) (get_global_id(0) + x, get_global_id(1) + y); \
+ float4 sourcePixel = read_imagef(source, sampler, sourceCoord); \
+ float4 destinationPixel = (float4) (__VA_ARGS__); \
+ write_imagef(destination, (int2) (get_global_id(0), get_global_id(1)), destinationPixel);
+
+static const char* colorMatrixKernelProgram =
+PROGRAM(
+const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
+
+__kernel void matrix(__write_only image2d_t destination, __read_only image2d_t source, float x, float y, __constant float *values)
+{
+ COLOR_MATRIX_KERNEL(values[0] * sourcePixel.x + values[1] * sourcePixel.y + values[2] * sourcePixel.z + values[3] * sourcePixel.w + values[4],
+ values[5] * sourcePixel.x + values[6] * sourcePixel.y + values[7] * sourcePixel.z + values[8] * sourcePixel.w + values[9],
+ values[10] * sourcePixel.x + values[11] * sourcePixel.y + values[12] * sourcePixel.z + values[13] * sourcePixel.w + values[14],
+ values[15] * sourcePixel.x + values[16] * sourcePixel.y + values[17] * sourcePixel.z + values[18] * sourcePixel.w + values[19])
+}
+
+__kernel void saturateAndHueRotate(__write_only image2d_t destination, __read_only image2d_t source, float x, float y, __constant float *components)
+{
+ COLOR_MATRIX_KERNEL(sourcePixel.x * components[0] + sourcePixel.y * components[1] + sourcePixel.z * components[2],
+ sourcePixel.x * components[3] + sourcePixel.y * components[4] + sourcePixel.z * components[5],
+ sourcePixel.x * components[6] + sourcePixel.y * components[7] + sourcePixel.z * components[8],
+ sourcePixel.w)
+}
+
+__kernel void luminance(__write_only image2d_t destination, __read_only image2d_t source, float x, float y)
+{
+ COLOR_MATRIX_KERNEL(0, 0, 0, 0.2125 * sourcePixel.x + 0.7154 * sourcePixel.y + 0.0721 * sourcePixel.z)
+}
+); // End of OpenCL kernels
+
+inline bool FilterContextOpenCL::compileFEColorMatrix()
+{
+ if (m_colorMatrixCompileStatus != openclNotCompiledYet)
+ return m_colorMatrixCompileStatus == openclCompileSuccessful;
+
+ m_colorMatrixCompileStatus = openclCompileFailed;
+ m_colorMatrixProgram = compileProgram(colorMatrixKernelProgram);
+ if (!m_colorMatrixProgram)
+ return false;
+
+ m_matrixOperation = kernelByName(m_colorMatrixProgram, "matrix");
+ if (!m_matrixOperation)
+ return false;
+ m_saturateAndHueRotateOperation = kernelByName(m_colorMatrixProgram, "saturateAndHueRotate");
+ if (!m_saturateAndHueRotateOperation)
+ return false;
+ m_luminanceOperation = kernelByName(m_colorMatrixProgram, "luminance");
+ if (!m_luminanceOperation)
+ return false;
+
+ m_colorMatrixCompileStatus = openclCompileSuccessful;
+ return openclCompileSuccessful;
+}
+
+inline void FilterContextOpenCL::applyFEColorMatrix(OpenCLHandle destination, IntSize destinationSize, OpenCLHandle source, IntPoint relativeSourceLocation, void* values, int type)
+{
+ cl_kernel colorMatrix;
+ OpenCLHandle clValues;
+
+ switch (type) {
+ case FECOLORMATRIX_TYPE_MATRIX:
+ colorMatrix = m_matrixOperation;
+ break;
+ case FECOLORMATRIX_TYPE_SATURATE:
+ colorMatrix = m_saturateAndHueRotateOperation;
+ break;
+ case FECOLORMATRIX_TYPE_HUEROTATE:
+ colorMatrix = m_saturateAndHueRotateOperation;
+ break;
+ case FECOLORMATRIX_TYPE_LUMINANCETOALPHA:
+ colorMatrix = m_luminanceOperation;
+ break;
+ default:
+ ASSERT_NOT_REACHED();
+ return;
+ }
+
+ RunKernel kernel(this, colorMatrix, destinationSize.width(), destinationSize.height());
+ kernel.addArgument(destination);
+ kernel.addArgument(source);
+ kernel.addArgument(relativeSourceLocation.x());
+ kernel.addArgument(relativeSourceLocation.y());
+ if (type == FECOLORMATRIX_TYPE_MATRIX)
+ clValues = kernel.addArgument(values, sizeof(float) * 20);
+ else if (type == FECOLORMATRIX_TYPE_SATURATE || type == FECOLORMATRIX_TYPE_HUEROTATE)
+ clValues = kernel.addArgument(values, sizeof(float) * 9);
+ kernel.run();
+
+ clValues.clear();
+}
+
+bool FEColorMatrix::platformApplyOpenCL()
+{
+ FilterContextOpenCL* context = FilterContextOpenCL::context();
+ if (!context || !context->compileFEColorMatrix())
+ return false;
+
+ FilterEffect* in = inputEffect(0);
+ OpenCLHandle source = in->openCLImage();
+ OpenCLHandle destination = createOpenCLImageResult();
+
+ IntPoint relativeSourceLocation(
+ absolutePaintRect().x() - in->absolutePaintRect().location().x(),
+ absolutePaintRect().y() - in->absolutePaintRect().location().y());
+
+ float components[9];
+ if (FECOLORMATRIX_TYPE_SATURATE == m_type)
+ calculateSaturateComponents(components, m_values[0]);
+ else if (FECOLORMATRIX_TYPE_HUEROTATE == m_type)
+ calculateHueRotateComponents(components, m_values[0]);
+
+ context->applyFEColorMatrix(destination, absolutePaintRect().size(), source, relativeSourceLocation, (FECOLORMATRIX_TYPE_MATRIX == m_type) ? m_values.data() : components, m_type);
+
+ return true;
+}
+
+} // namespace WebCore
+
+#endif // ENABLE(FILTERS) && ENABLE(OPENCL)
