1/*
2 * Copyright (C) 2013 University of Szeged
3 * Copyright (C) 2013 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
28#include "config.h"
29#if ENABLE(FILTERS) && ENABLE(OPENCL)
30#include "FEMerge.h"
31
32#include "FilterContextOpenCL.h"
33
34namespace WebCore {
35
36static const char* mergeKernelProgram =
37PROGRAM(
38const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
39__kernel void copy(__write_only image2d_t destination, __read_only image2d_t source, int x, int y)
40{
41 float4 destinationPixel = read_imagef(source, sampler, (int2) (get_global_id(0) + x, get_global_id(1) + y));
42 write_imagef(destination, (int2) (get_global_id(0), get_global_id(1)), destinationPixel);
43}
44
45__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)
46{
47 int2 destinationCoord = (int2) (get_global_id(0), get_global_id(1));
48 int2 sourceCoordA = (int2) (destinationCoord.x + xA, destinationCoord.y + yA);
49 int2 sourceCoordB = (int2) (destinationCoord.x + xB, destinationCoord.y + yB);
50 float4 destinationPixel = read_imagef(previousDestination, sampler, destinationCoord);
51 float4 sourcePixelA = read_imagef(sourceA, sampler, sourceCoordA);
52 float4 sourcePixelB = read_imagef(sourceB, sampler, sourceCoordB);
53
54 destinationPixel = sourcePixelA + destinationPixel * (1 - sourcePixelA.w);
55 destinationPixel = sourcePixelB + destinationPixel * (1 - sourcePixelB.w);
56
57 write_imagef(destination, destinationCoord, destinationPixel);
58}
59); // End of OpenCL kernels
60
61inline bool FilterContextOpenCL::compileFEMerge()
62{
63 if (m_mergeWasCompiled || inError())
64 return !inError();
65
66 m_mergeWasCompiled = true;
67
68 if (isResourceAllocationFailed((m_mergeProgram = compileProgram(mergeKernelProgram))))
69 return false;
70 if (isResourceAllocationFailed((m_mergeCopyOperation = kernelByName(m_mergeProgram, "copy"))))
71 return false;
72 if (isResourceAllocationFailed((m_mergeOperation = kernelByName(m_mergeProgram, "merge"))))
73 return false;
74 return true;
75}
76
77inline void FilterContextOpenCL::applyFEMergeCopy(OpenCLHandle destination, IntSize destinationSize, OpenCLHandle source, IntPoint &relativeSourcePoint)
78{
79 RunKernel kernel(this, m_mergeCopyOperation, destinationSize.width(), destinationSize.height());
80 kernel.addArgument(destination);
81 kernel.addArgument(source);
82 kernel.addArgument(relativeSourcePoint.x());
83 kernel.addArgument(relativeSourcePoint.y());
84 kernel.run();
85}
86
87inline void FilterContextOpenCL::applyFEMerge(OpenCLHandle destination, OpenCLHandle previousDestination, OpenCLHandle sourceA, OpenCLHandle sourceB, IntSize destinationSize, IntPoint &relativeSourcePointA, IntPoint &relativeSourcePointB)
88{
89 RunKernel kernel(this, m_mergeOperation, destinationSize.width(), destinationSize.height());
90 kernel.addArgument(destination);
91 kernel.addArgument(previousDestination);
92 kernel.addArgument(sourceA);
93 kernel.addArgument(sourceB);
94 kernel.addArgument(relativeSourcePointA.x());
95 kernel.addArgument(relativeSourcePointA.y());
96 kernel.addArgument(relativeSourcePointB.x());
97 kernel.addArgument(relativeSourcePointB.y());
98 kernel.run();
99}
100
101bool FEMerge::platformApplyOpenCL()
102{
103 FilterContextOpenCL* context = FilterContextOpenCL::context();
104 if (!context)
105 return false;
106
107 context->compileFEMerge();
108
109 unsigned size = numberOfEffectInputs();
110 ASSERT(size > 0);
111
112 OpenCLHandle destination = createOpenCLImageResult();
113 OpenCLHandle sourceA = 0;
114 OpenCLHandle sourceB = 0;
115 FilterEffect* in;
116
117 int i = 0;
118
119 if (size & 1) {
120 in = inputEffect(i++);
121 sourceA = in->openCLImage();
122 IntPoint relativeSourcePoint(in->absolutePaintRect().location());
123 relativeSourcePoint.setX(absolutePaintRect().x() - relativeSourcePoint.x());
124 relativeSourcePoint.setY(absolutePaintRect().y() - relativeSourcePoint.y());
125
126 context->applyFEMergeCopy(destination, absolutePaintRect().size(), sourceA, relativeSourcePoint);
127 if (size == 1)
128 return true;
129 } else
130 context->fill(destination, absolutePaintRect().size(), Color(0.0f, 0.0f, 0.0f, 0.0f));
131
132 OpenCLHandle previousDestination = context->createOpenCLImage(absolutePaintRect().size());
133
134 while (i < size) {
135 OpenCLHandle temp = previousDestination;
136 previousDestination = destination;
137 destination = temp;
138
139 in = inputEffect(i++);
140 sourceA = in->openCLImage();
141 IntPoint relativeSourcePointA(in->absolutePaintRect().location());
142 relativeSourcePointA.setX(absolutePaintRect().x() - relativeSourcePointA.x());
143 relativeSourcePointA.setY(absolutePaintRect().y() - relativeSourcePointA.y());
144
145 in = inputEffect(i++);
146 sourceB = in->openCLImage();
147 IntPoint relativeSourcePointB(in->absolutePaintRect().location());
148 relativeSourcePointB.setX(absolutePaintRect().x() - relativeSourcePointB.x());
149 relativeSourcePointB.setY(absolutePaintRect().y() - relativeSourcePointB.y());
150
151 context->applyFEMerge(destination, previousDestination, sourceA, sourceB, absolutePaintRect().size(), relativeSourcePointA, relativeSourcePointB);
152 }
153 setOpenCLImage(destination);
154 previousDestination.clear();
155 return true;
156}
157
158} // namespace WebCore
159
160#endif // ENABLE(FILTERS) && ENABLE(OPENCL)