Modified: trunk/Source/WebCore/platform/graphics/gpu/opencl/FilterContextOpenCL.h (143577 => 143578)
--- trunk/Source/WebCore/platform/graphics/gpu/opencl/FilterContextOpenCL.h 2013-02-21 10:26:48 UTC (rev 143577)
+++ trunk/Source/WebCore/platform/graphics/gpu/opencl/FilterContextOpenCL.h 2013-02-21 10:43:03 UTC (rev 143578)
@@ -54,6 +54,10 @@
, m_fillWasCompiled(false)
, m_fillProgram(0)
, m_fill(0)
+ , m_mergeWasCompiled(false)
+ , m_mergeProgram(0)
+ , m_mergeCopyOperation(0)
+ , m_mergeOperation(0)
, m_colorMatrixWasCompiled(false)
, m_colorMatrixProgram(0)
, m_matrixOperation(0)
@@ -90,7 +94,10 @@
inline bool compileFEColorMatrix();
inline bool compileFETurbulence();
+ inline bool compileFEMerge();
+ inline void applyFEMergeCopy(OpenCLHandle, IntSize, OpenCLHandle, IntPoint&);
+ inline void applyFEMerge(OpenCLHandle, OpenCLHandle, OpenCLHandle, OpenCLHandle, IntSize, IntPoint&, IntPoint&);
inline void applyFEColorMatrix(OpenCLHandle, IntSize, OpenCLHandle, IntPoint, float*, int);
inline void applyFETurbulence(OpenCLHandle, IntSize, int, void*, void*, void*, void*, void*,
int*, int, int, int, int, float, float, bool, int, int);
@@ -187,10 +194,15 @@
cl_program m_fillProgram;
cl_kernel m_fill;
+ bool m_mergeWasCompiled;
+ cl_program m_mergeProgram;
+ cl_kernel m_mergeCopyOperation;
+ cl_kernel m_mergeOperation;
+
bool m_colorMatrixWasCompiled;
cl_program m_colorMatrixProgram;
cl_kernel m_matrixOperation;
- cl_kernel m_saturateAndHueRotateOperation;
+ cl_kernel m_saturateAndHueRotateOperation;
cl_kernel m_luminanceOperation;
bool m_turbulenceWasCompiled;
Added: trunk/Source/WebCore/platform/graphics/gpu/opencl/OpenCLFEMerge.cpp (0 => 143578)
--- trunk/Source/WebCore/platform/graphics/gpu/opencl/OpenCLFEMerge.cpp (rev 0)
+++ trunk/Source/WebCore/platform/graphics/gpu/opencl/OpenCLFEMerge.cpp 2013-02-21 10:43:03 UTC (rev 143578)
@@ -0,0 +1,161 @@
+/*
+ * 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 "FEMerge.h"
+
+#include "FilterContextOpenCL.h"
+
+namespace WebCore {
+
+static const char* mergeKernelProgram =
+PROGRAM(
+const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
+
+__kernel void copy(__write_only image2d_t destination, __read_only image2d_t source, int x, int y)
+{
+ float4 destinationPixel = read_imagef(source, sampler, (int2) (get_global_id(0) + x, get_global_id(1) + y));
+ write_imagef(destination, (int2) (get_global_id(0), get_global_id(1)), destinationPixel);
+}
+
+__kernel void merge(__write_only image2d_t destination, __read_only image2d_t previousDestination, __read_only image2d_t sourceA, __read_only image2d_t sourceB, int xA, int yA, int xB, int yB)
+{
+ int2 destinationCoord = (int2) (get_global_id(0), get_global_id(1));
+ int2 sourceCoordA = (int2) (destinationCoord.x + xA, destinationCoord.y + yA);
+ int2 sourceCoordB = (int2) (destinationCoord.x + xB, destinationCoord.y + yB);
+ float4 destinationPixel = read_imagef(previousDestination, sampler, destinationCoord);
+ float4 sourcePixelA = read_imagef(sourceA, sampler, sourceCoordA);
+ float4 sourcePixelB = read_imagef(sourceB, sampler, sourceCoordB);
+
+ destinationPixel = sourcePixelA + destinationPixel * (1 - sourcePixelA.w);
+ destinationPixel = sourcePixelB + destinationPixel * (1 - sourcePixelB.w);
+
+ write_imagef(destination, destinationCoord, destinationPixel);
+}
+); // End of OpenCL kernels
+
+inline bool FilterContextOpenCL::compileFEMerge()
+{
+ if (m_mergeWasCompiled || inError())
+ return !inError();
+
+ m_mergeWasCompiled = true;
+
+ if (isResourceAllocationFailed((m_mergeProgram = compileProgram(mergeKernelProgram))))
+ return false;
+ if (isResourceAllocationFailed((m_mergeCopyOperation = kernelByName(m_mergeProgram, "copy"))))
+ return false;
+ if (isResourceAllocationFailed((m_mergeOperation = kernelByName(m_mergeProgram, "merge"))))
+ return false;
+ return true;
+}
+
+inline void FilterContextOpenCL::applyFEMergeCopy(OpenCLHandle destination, IntSize destinationSize, OpenCLHandle source, IntPoint& relativeSourcePoint)
+{
+ RunKernel kernel(this, m_mergeCopyOperation, destinationSize.width(), destinationSize.height());
+ kernel.addArgument(destination);
+ kernel.addArgument(source);
+ kernel.addArgument(relativeSourcePoint.x());
+ kernel.addArgument(relativeSourcePoint.y());
+ kernel.run();
+}
+
+inline void FilterContextOpenCL::applyFEMerge(OpenCLHandle destination, OpenCLHandle previousDestination, OpenCLHandle sourceA, OpenCLHandle sourceB, IntSize destinationSize, IntPoint& relativeSourcePointA, IntPoint& relativeSourcePointB)
+{
+ RunKernel kernel(this, m_mergeOperation, destinationSize.width(), destinationSize.height());
+ kernel.addArgument(destination);
+ kernel.addArgument(previousDestination);
+ kernel.addArgument(sourceA);
+ kernel.addArgument(sourceB);
+ kernel.addArgument(relativeSourcePointA.x());
+ kernel.addArgument(relativeSourcePointA.y());
+ kernel.addArgument(relativeSourcePointB.x());
+ kernel.addArgument(relativeSourcePointB.y());
+ kernel.run();
+}
+
+bool FEMerge::platformApplyOpenCL()
+{
+ FilterContextOpenCL* context = FilterContextOpenCL::context();
+ if (!context)
+ return false;
+
+ context->compileFEMerge();
+
+ unsigned size = numberOfEffectInputs();
+ ASSERT(size > 0);
+
+ OpenCLHandle destination = createOpenCLImageResult();
+ OpenCLHandle sourceA = 0;
+ OpenCLHandle sourceB = 0;
+ FilterEffect* in;
+
+ int i = 0;
+
+ if (size & 1) {
+ in = inputEffect(i++);
+ sourceA = in->openCLImage();
+ IntPoint relativeSourcePoint(in->absolutePaintRect().location());
+ relativeSourcePoint.setX(absolutePaintRect().x() - relativeSourcePoint.x());
+ relativeSourcePoint.setY(absolutePaintRect().y() - relativeSourcePoint.y());
+
+ context->applyFEMergeCopy(destination, absolutePaintRect().size(), sourceA, relativeSourcePoint);
+ if (size == 1)
+ return true;
+ } else
+ context->fill(destination, absolutePaintRect().size(), Color(0.0f, 0.0f, 0.0f, 0.0f));
+
+ OpenCLHandle previousDestination = context->createOpenCLImage(absolutePaintRect().size());
+
+ while (i < size) {
+ OpenCLHandle temp = previousDestination;
+ previousDestination = destination;
+ destination = temp;
+
+ in = inputEffect(i++);
+ sourceA = in->openCLImage();
+ IntPoint relativeSourcePointA(in->absolutePaintRect().location());
+ relativeSourcePointA.setX(absolutePaintRect().x() - relativeSourcePointA.x());
+ relativeSourcePointA.setY(absolutePaintRect().y() - relativeSourcePointA.y());
+
+ in = inputEffect(i++);
+ sourceB = in->openCLImage();
+ IntPoint relativeSourcePointB(in->absolutePaintRect().location());
+ relativeSourcePointB.setX(absolutePaintRect().x() - relativeSourcePointB.x());
+ relativeSourcePointB.setY(absolutePaintRect().y() - relativeSourcePointB.y());
+
+ context->applyFEMerge(destination, previousDestination, sourceA, sourceB, absolutePaintRect().size(), relativeSourcePointA, relativeSourcePointB);
+ }
+ setOpenCLImage(destination);
+ previousDestination.clear();
+ return true;
+}
+
+} // namespace WebCore
+
+#endif // ENABLE(FILTERS) && ENABLE(OPENCL)
