Diff
Modified: trunk/Source/WebCore/ChangeLog (142734 => 142735)
--- trunk/Source/WebCore/ChangeLog 2013-02-13 11:19:08 UTC (rev 142734)
+++ trunk/Source/WebCore/ChangeLog 2013-02-13 11:39:32 UTC (rev 142735)
@@ -1,3 +1,25 @@
+2013-02-13 Tamas Czene <[email protected]>
+
+ OpenCL implementation of Flood SVG filters.
+ https://bugs.webkit.org/show_bug.cgi?id=109580
+
+ Reviewed by Zoltan Herczeg.
+
+ * Target.pri:
+ * platform/graphics/filters/FEFlood.h:
+ (FEFlood):
+ * platform/graphics/gpu/opencl/FilterContextOpenCL.cpp:
+ (WebCore):
+ (WebCore::PROGRAM_STR):
+ (WebCore::FilterContextOpenCL::compileFill):
+ (WebCore::FilterContextOpenCL::fill):
+ * platform/graphics/gpu/opencl/FilterContextOpenCL.h:
+ (WebCore::FilterContextOpenCL::FilterContextOpenCL):
+ (FilterContextOpenCL):
+ * platform/graphics/gpu/opencl/OpenCLFEFlood.cpp: Added.
+ (WebCore):
+ (WebCore::FEFlood::platformApplyOpenCL):
+
2013-02-13 Mike West <[email protected]>
location.href does not throw SECURITY_ERR when accessed across origins with JSC bindings
Modified: trunk/Source/WebCore/Target.pri (142734 => 142735)
--- trunk/Source/WebCore/Target.pri 2013-02-13 11:19:08 UTC (rev 142734)
+++ trunk/Source/WebCore/Target.pri 2013-02-13 11:39:32 UTC (rev 142735)
@@ -4219,6 +4219,7 @@
SOURCES += \
platform/graphics/gpu/opencl/FilterContextOpenCL.cpp \
platform/graphics/gpu/opencl/OpenCLFEColorMatrix.cpp \
+ platform/graphics/gpu/opencl/OpenCLFEFlood.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/FEFlood.h (142734 => 142735)
--- trunk/Source/WebCore/platform/graphics/filters/FEFlood.h 2013-02-13 11:19:08 UTC (rev 142734)
+++ trunk/Source/WebCore/platform/graphics/filters/FEFlood.h 2013-02-13 11:39:32 UTC (rev 142735)
@@ -40,6 +40,9 @@
bool setFloodOpacity(float);
virtual void platformApplySoftware();
+#if ENABLE(OPENCL)
+ virtual bool platformApplyOpenCL();
+#endif
virtual void dump();
virtual void determineAbsolutePaintRect() { setAbsolutePaintRect(enclosingIntRect(maxEffectRect())); }
Modified: trunk/Source/WebCore/platform/graphics/gpu/opencl/FilterContextOpenCL.cpp (142734 => 142735)
--- trunk/Source/WebCore/platform/graphics/gpu/opencl/FilterContextOpenCL.cpp 2013-02-13 11:19:08 UTC (rev 142734)
+++ trunk/Source/WebCore/platform/graphics/gpu/opencl/FilterContextOpenCL.cpp 2013-02-13 11:39:32 UTC (rev 142735)
@@ -225,6 +225,49 @@
source = destination;
}
+static const char* fillKernelProgram =
+PROGRAM_STR(
+__kernel void fill(__write_only image2d_t destination, float r, float g, float b, float a)
+{
+ float4 sourcePixel = (float4)(r, g, b, a);
+ write_imagef(destination, (int2)(get_global_id(0), get_global_id(1)), sourcePixel);
+}
+);
+
+inline bool FilterContextOpenCL::compileFill()
+{
+ if (m_fillWasCompiled || inError())
+ return !inError();
+
+ m_fillWasCompiled = true;
+
+ if (isResourceAllocationFailed((m_fillProgram = compileProgram(fillKernelProgram))))
+ return false;
+ if (isResourceAllocationFailed((m_fill = kernelByName(m_fillProgram, "fill"))))
+ return false;
+ return true;
+}
+
+void FilterContextOpenCL::fill(cl_mem image, IntSize imageSize, Color color)
+{
+ if (!m_context || inError())
+ return;
+
+ compileFill();
+
+ float r, g, b, a;
+
+ color.getRGBA(r, g, b, a);
+
+ RunKernel kernel(this, m_fill, imageSize.width(), imageSize.height());
+ kernel.addArgument(image);
+ kernel.addArgument(r);
+ kernel.addArgument(g);
+ kernel.addArgument(b);
+ kernel.addArgument(a);
+ kernel.run();
+}
+
cl_program FilterContextOpenCL::compileProgram(const char* source)
{
cl_program program;
Modified: trunk/Source/WebCore/platform/graphics/gpu/opencl/FilterContextOpenCL.h (142734 => 142735)
--- trunk/Source/WebCore/platform/graphics/gpu/opencl/FilterContextOpenCL.h 2013-02-13 11:19:08 UTC (rev 142734)
+++ trunk/Source/WebCore/platform/graphics/gpu/opencl/FilterContextOpenCL.h 2013-02-13 11:39:32 UTC (rev 142735)
@@ -51,6 +51,9 @@
, m_transformColorSpaceWasCompiled(false)
, m_transformColorSpaceProgram(0)
, m_transformColorSpaceKernel(0)
+ , m_fillWasCompiled(false)
+ , m_fillProgram(0)
+ , m_fill(0)
, m_colorMatrixWasCompiled(false)
, m_colorMatrixProgram(0)
, m_matrixOperation(0)
@@ -79,6 +82,9 @@
OpenCLHandle createOpenCLImage(IntSize);
+ inline bool compileFill();
+ void fill(cl_mem, IntSize, Color);
+
inline bool compileTransformColorSpaceProgram();
void openCLTransformColorSpace(OpenCLHandle&, IntRect, ColorSpace, ColorSpace);
@@ -177,6 +183,10 @@
cl_program m_transformColorSpaceProgram;
cl_kernel m_transformColorSpaceKernel;
+ bool m_fillWasCompiled;
+ cl_program m_fillProgram;
+ cl_kernel m_fill;
+
bool m_colorMatrixWasCompiled;
cl_program m_colorMatrixProgram;
cl_kernel m_matrixOperation;
Added: trunk/Source/WebCore/platform/graphics/gpu/opencl/OpenCLFEFlood.cpp (0 => 142735)
--- trunk/Source/WebCore/platform/graphics/gpu/opencl/OpenCLFEFlood.cpp (rev 0)
+++ trunk/Source/WebCore/platform/graphics/gpu/opencl/OpenCLFEFlood.cpp 2013-02-13 11:39:32 UTC (rev 142735)
@@ -0,0 +1,57 @@
+/*
+ * Copyright (C) 2013 University of Szeged
+ * Copyright (C) 2013 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 "FEFlood.h"
+
+#include "FilterContextOpenCL.h"
+
+namespace WebCore {
+
+bool FEFlood::platformApplyOpenCL()
+{
+ FilterContextOpenCL* context = FilterContextOpenCL::context();
+
+ if (!context)
+ return false;
+
+ if (context->inError())
+ return true;
+
+ cl_mem destination = createOpenCLImageResult();
+
+ Color color = colorWithOverrideAlpha(floodColor().rgb(), floodOpacity());
+
+ context->fill(destination, absolutePaintRect().size(), color);
+
+ return true;
+}
+
+} // namespace WebCore
+
+#endif // ENABLE(FILTERS) && ENABLE(OPENCL)
