OpenCL implementation of Flood SVG filters.
authorcommit-queue@webkit.org <commit-queue@webkit.org@268f45cc-cd09-0410-ab3c-d52691b4dbfc>
Wed, 13 Feb 2013 11:39:32 +0000 (11:39 +0000)
committercommit-queue@webkit.org <commit-queue@webkit.org@268f45cc-cd09-0410-ab3c-d52691b4dbfc>
Wed, 13 Feb 2013 11:39:32 +0000 (11:39 +0000)
https://bugs.webkit.org/show_bug.cgi?id=109580

Patch by Tamas Czene <tczene@inf.u-szeged.hu> on 2013-02-13
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):

git-svn-id: https://svn.webkit.org/repository/webkit/trunk@142735 268f45cc-cd09-0410-ab3c-d52691b4dbfc

Source/WebCore/ChangeLog
Source/WebCore/Target.pri
Source/WebCore/platform/graphics/filters/FEFlood.h
Source/WebCore/platform/graphics/gpu/opencl/FilterContextOpenCL.cpp
Source/WebCore/platform/graphics/gpu/opencl/FilterContextOpenCL.h
Source/WebCore/platform/graphics/gpu/opencl/OpenCLFEFlood.cpp [new file with mode: 0644]

index 0d69b93d3311e424eca2c40f459b93ad60eb36c5..3a9192f8d7546f37ca128cf3eb685d0f01013a07 100644 (file)
@@ -1,3 +1,25 @@
+2013-02-13  Tamas Czene  <tczene@inf.u-szeged.hu>
+
+        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  <mkwst@chromium.org>
 
         location.href does not throw SECURITY_ERR when accessed across origins with JSC bindings
index fcfca9fa232633240693616d48ed63fee819cab5..fb1862eb2b6e144e323c2edb2bbbd718d7318b7c 100644 (file)
@@ -4219,6 +4219,7 @@ contains(DEFINES, ENABLE_OPENCL=1) {
     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
index e0484b44929d202729ece5770b0651b3da94716f..e78b6450229bf6807ea293766c89b933b3beb4d7 100644 (file)
@@ -40,6 +40,9 @@ public:
     bool setFloodOpacity(float);
 
     virtual void platformApplySoftware();
+#if ENABLE(OPENCL)
+    virtual bool platformApplyOpenCL();
+#endif
     virtual void dump();
 
     virtual void determineAbsolutePaintRect() { setAbsolutePaintRect(enclosingIntRect(maxEffectRect())); }
index 580933d21dab73a569d6f36e5377048df404d88b..72f58de6aa5a40add6ceadfcddf81c8aa1e02e97 100644 (file)
@@ -225,6 +225,49 @@ void FilterContextOpenCL::openCLTransformColorSpace(OpenCLHandle& source, IntRec
     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;
index 80264208379ae74019f6865f193f21f900134420..9a71a55b3a2ee98070919a22502ee111d7c5d4b5 100644 (file)
@@ -51,6 +51,9 @@ public:
         , 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 @@ public:
 
     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 @@ private:
     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;
diff --git a/Source/WebCore/platform/graphics/gpu/opencl/OpenCLFEFlood.cpp b/Source/WebCore/platform/graphics/gpu/opencl/OpenCLFEFlood.cpp
new file mode 100644 (file)
index 0000000..6ef1f66
--- /dev/null
@@ -0,0 +1,57 @@
+/*
+ * Copyright (C) 2013 University of Szeged
+ * Copyright (C) 2013 Tamas Czene <tczene@inf.u-szeged.hu>
+ * 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)