Attachment #163638 for bug #70099

View Details Raw Unified Return to Bug 70099

| Differences between



    DEFINES += QT_OPENGL_SHIMS=1
}

contains(DEFINES, ENABLE_OPENCL=1) {
    HEADERS += platform/graphics/filters/OpenCL/OpenCLContext.h
    SOURCES += platform/graphics/filters/OpenCL/OpenCLContext.cpp
    SOURCES += platform/graphics/filters/OpenCL/OpenCLFEImage.cpp
    SOURCES += platform/graphics/filters/OpenCL/OpenCLFEComposite.cpp
    SOURCES += platform/graphics/filters/OpenCL/OpenCLFEBlend.cpp
    SOURCES += platform/graphics/filters/OpenCL/OpenCLFEFlood.cpp
    SOURCES += platform/graphics/filters/OpenCL/OpenCLFEOffset.cpp
    SOURCES += platform/graphics/filters/OpenCL/OpenCLFESourceGraphic.cpp
    SOURCES += platform/graphics/filters/OpenCL/OpenCLFESourceAlpha.cpp
    SOURCES += platform/graphics/filters/OpenCL/OpenCLFETurbulence.cpp
    SOURCES += platform/graphics/filters/OpenCL/OpenCLFEColorMatrix.cpp
    SOURCES += platform/graphics/filters/OpenCL/OpenCLFEConvolveMatrix.cpp
    #SOURCES += platform/graphics/filters/OpenCL/OpenCLFEComponentTransfer.cpp
    SOURCES += platform/graphics/filters/OpenCL/OpenCLFEMerge.cpp
    SOURCES += platform/graphics/filters/OpenCL/OpenCLFEMorphology.cpp
    SOURCES += platform/graphics/filters/OpenCL/OpenCLFETile.cpp
    DEFINES += OPENCL=1
}

use?(GRAPHICS_SURFACE) {
    mac {
        SOURCES += platform/graphics/surfaces/mac/GraphicsSurfaceMac.cpp

Lines 263-268 unix\|win32-g++* { a/Source/WebCore/WebCore.pri_sec1
- a/Source/WebCore/WebCore.pri +5 lines
263	QMAKE_PKGCONFIG_REQUIRES = QtCore QtGui QtNetwork QtWidgets	263	QMAKE_PKGCONFIG_REQUIRES = QtCore QtGui QtNetwork QtWidgets
264	}	264	}
265		265
		266	contains(DEFINES, ENABLE_OPENCL=1){
		267	QMAKE_LFLAGS += -I/usr/include/ -L/usr/lib/
		268	LIBS += -lOpenCL
		269	}
		270
266	# Disable C++0x mode in WebCore for those who enabled it in their Qt's mkspec	271	# Disable C++0x mode in WebCore for those who enabled it in their Qt's mkspec
267	-g++:QMAKE_CXXFLAGS -= -std=c++0x -std=gnu++0x	272	-g++:QMAKE_CXXFLAGS -= -std=c++0x -std=gnu++0x
268		273

Lines 52-58 public: a/Source/WebCore/platform/graphics/filters/FEBlend.h_sec1
- a/Source/WebCore/platform/graphics/filters/FEBlend.h -1 / +3 lines
52	#if USE(SKIA)	52	#if USE(SKIA)
53	virtual bool platformApplySkia();	53	virtual bool platformApplySkia();
54	#endif	54	#endif
55		55	#if ENABLE(OPENCL)
		56	virtual bool platformApplyOpenCL();
		57	#endif
56	virtual void platformApplySoftware();	58	virtual void platformApplySoftware();
57	virtual void dump();	59	virtual void dump();
58		60

Lines 49-54 public: a/Source/WebCore/platform/graphics/filters/FEColorMatrix.h_sec1
- a/Source/WebCore/platform/graphics/filters/FEColorMatrix.h +3 lines
49	bool setValues(const Vector<float>&);	49	bool setValues(const Vector<float>&);
50		50
51	virtual void platformApplySoftware();	51	virtual void platformApplySoftware();
		52	#if ENABLE(OPENCL)
		53	virtual bool platformApplyOpenCL();
		54	#endif
52	#if USE(SKIA)	55	#if USE(SKIA)
53	virtual bool platformApplySkia();	56	virtual bool platformApplySkia();
54	#endif	57	#endif



    }
}

#if ENABLE(OPENCL)
bool FEComponentTransfer::platformApplyOpenCL()
{
    asImageBuffer();
    platformApplySoftware();
    ImageBuffer* sourceImage = asImageBuffer();
    RefPtr<Uint8ClampedArray> sourceImageData = sourceImage->getUnmultipliedImageData(IntRect(IntPoint(),sourceImage->internalSize()));
    createOpenCLImageResult(sourceImageData->data());
    return true;
}
#endif

void FEComponentTransfer::getValues(unsigned char rValues[256], unsigned char gValues[256], unsigned char bValues[256], unsigned char aValues[256])
{
    for (unsigned i = 0; i < 256; ++i)

Lines 79-84 public: a/Source/WebCore/platform/graphics/filters/FEComponentTransfer.h_sec1
- a/Source/WebCore/platform/graphics/filters/FEComponentTransfer.h +3 lines
79	void setAlphaFunction(const ComponentTransferFunction&);	79	void setAlphaFunction(const ComponentTransferFunction&);
80		80
81	virtual void platformApplySoftware();	81	virtual void platformApplySoftware();
		82	#if ENABLE(OPENCL)
		83	virtual bool platformApplyOpenCL();
		84	#endif
82	#if USE(SKIA)	85	#if USE(SKIA)
83	virtual bool platformApplySkia();	86	virtual bool platformApplySkia();
84	#endif	87	#endif

Lines 62-67 public: a/Source/WebCore/platform/graphics/filters/FEComposite.h_sec1
- a/Source/WebCore/platform/graphics/filters/FEComposite.h +3 lines
62	virtual void correctFilterResultIfNeeded() OVERRIDE;	62	virtual void correctFilterResultIfNeeded() OVERRIDE;
63		63
64	virtual void platformApplySoftware();	64	virtual void platformApplySoftware();
		65	#if ENABLE(OPENCL)
		66	virtual bool platformApplyOpenCL();
		67	#endif
65	virtual void dump();	68	virtual void dump();
66		69
67	virtual void determineAbsolutePaintRect();	70	virtual void determineAbsolutePaintRect();

Lines 71-76 public: a/Source/WebCore/platform/graphics/filters/FEConvolveMatrix.h_sec1
- a/Source/WebCore/platform/graphics/filters/FEConvolveMatrix.h +3 lines
71	bool setPreserveAlpha(bool);	71	bool setPreserveAlpha(bool);
72		72
73	virtual void platformApplySoftware();	73	virtual void platformApplySoftware();
		74	#if ENABLE(OPENCL)
		75	virtual bool platformApplyOpenCL();
		76	#endif
74	virtual void dump();	77	virtual void dump();
75		78
76	virtual void determineAbsolutePaintRect() { setAbsolutePaintRect(enclosingIntRect(maxEffectRect())); }	79	virtual void determineAbsolutePaintRect() { setAbsolutePaintRect(enclosingIntRect(maxEffectRect())); }



        clearShaderResult();
}

#if ENABLE(OPENCL)
bool FECustomFilter::platformApplyOpenCL()
{
    asImageBuffer();
    platformApplySoftware();
    ImageBuffer* sourceImage = asImageBuffer();
    RefPtr<Uint8ClampedArray> sourceImageData = sourceImage->getUnmultipliedImageData(IntRect(IntPoint(),sourceImage->internalSize()));
    createOpenCLImageResult(sourceImageData->data());
    return true;
}
#endif

void FECustomFilter::clearShaderResult()
{
    clearResult();

Lines 64-69 public: a/Source/WebCore/platform/graphics/filters/FECustomFilter.h_sec1
- a/Source/WebCore/platform/graphics/filters/FECustomFilter.h +3 lines
64	CustomFilterOperation::MeshType);	64	CustomFilterOperation::MeshType);
65		65
66	virtual void platformApplySoftware();	66	virtual void platformApplySoftware();
		67	#if ENABLE(OPENCL)
		68	virtual bool platformApplyOpenCL();
		69	#endif
67	virtual void dump();	70	virtual void dump();
68		71
69	virtual TextStream& externalRepresentation(TextStream&, int indention) const;	72	virtual TextStream& externalRepresentation(TextStream&, int indention) const;



    }
}

#if ENABLE(OPENCL)
bool FEDisplacementMap::platformApplyOpenCL()
{
    asImageBuffer();
    platformApplySoftware();
    ImageBuffer* sourceImage = asImageBuffer();
    RefPtr<Uint8ClampedArray> sourceImageData = sourceImage->getUnmultipliedImageData(IntRect(IntPoint(),sourceImage->internalSize()));
    createOpenCLImageResult(sourceImageData->data());
    return true;
}
#endif

void FEDisplacementMap::dump()
{
}

Lines 51-56 public: a/Source/WebCore/platform/graphics/filters/FEDisplacementMap.h_sec1
- a/Source/WebCore/platform/graphics/filters/FEDisplacementMap.h +3 lines
51	bool setScale(float);	51	bool setScale(float);
52		52
53	virtual void platformApplySoftware();	53	virtual void platformApplySoftware();
		54	#if ENABLE(OPENCL)
		55	virtual bool platformApplyOpenCL();
		56	#endif
54	virtual void dump();	57	virtual void dump();
55		58
56	virtual void determineAbsolutePaintRect() { setAbsolutePaintRect(enclosingIntRect(maxEffectRect())); }	59	virtual void determineAbsolutePaintRect() { setAbsolutePaintRect(enclosingIntRect(maxEffectRect())); }



    resultImage->context()->drawImageBuffer(sourceImage, ColorSpaceDeviceRGB, drawingRegion);
}

#if ENABLE(OPENCL)
bool FEDropShadow::platformApplyOpenCL()
{
    asImageBuffer();
    platformApplySoftware();
    ImageBuffer* sourceImage = asImageBuffer();
    RefPtr<Uint8ClampedArray> sourceImageData = sourceImage->getUnmultipliedImageData(IntRect(IntPoint(),sourceImage->internalSize()));
    createOpenCLImageResult(sourceImageData->data());
    return true;
}
#endif

void FEDropShadow::dump()
{
}

Lines 52-57 public: a/Source/WebCore/platform/graphics/filters/FEDropShadow.h_sec1
- a/Source/WebCore/platform/graphics/filters/FEDropShadow.h +3 lines
52	static float calculateStdDeviation(float);	52	static float calculateStdDeviation(float);
53		53
54	virtual void platformApplySoftware();	54	virtual void platformApplySoftware();
		55	#if ENABLE(OPENCL)
		56	virtual bool platformApplyOpenCL();
		57	#endif
55	virtual void dump();	58	virtual void dump();
56		59
57	virtual void determineAbsolutePaintRect();	60	virtual void determineAbsolutePaintRect();

Lines 40-45 public: a/Source/WebCore/platform/graphics/filters/FEFlood.h_sec1
- a/Source/WebCore/platform/graphics/filters/FEFlood.h +3 lines
40	bool setFloodOpacity(float);	40	bool setFloodOpacity(float);
41		41
42	virtual void platformApplySoftware();	42	virtual void platformApplySoftware();
		43	#if ENABLE(OPENCL)
		44	virtual bool platformApplyOpenCL();
		45	#endif
43	virtual void dump();	46	virtual void dump();
44		47
45	virtual void determineAbsolutePaintRect() { setAbsolutePaintRect(enclosingIntRect(maxEffectRect())); }	48	virtual void determineAbsolutePaintRect() { setAbsolutePaintRect(enclosingIntRect(maxEffectRect())); }



    platformApply(srcPixelArray, tmpPixelArray, kernelSizeX, kernelSizeY, paintSize);
}

#if ENABLE(OPENCL)
bool FEGaussianBlur::platformApplyOpenCL()
{
    asImageBuffer();
    platformApplySoftware();
    ImageBuffer* sourceImage = asImageBuffer();
    RefPtr<Uint8ClampedArray> sourceImageData = sourceImage->getUnmultipliedImageData(IntRect(IntPoint(),sourceImage->internalSize()));
    createOpenCLImageResult(sourceImageData->data());
    return true;
}
#endif

void FEGaussianBlur::dump()
{
}

Lines 41-46 public: a/Source/WebCore/platform/graphics/filters/FEGaussianBlur.h_sec1
- a/Source/WebCore/platform/graphics/filters/FEGaussianBlur.h +3 lines
41	static float calculateStdDeviation(float);	41	static float calculateStdDeviation(float);
42		42
43	virtual void platformApplySoftware();	43	virtual void platformApplySoftware();
		44	#if ENABLE(OPENCL)
		45	virtual bool platformApplyOpenCL();
		46	#endif
44	virtual void dump();	47	virtual void dump();
45		48
46	virtual void determineAbsolutePaintRect();	49	virtual void determineAbsolutePaintRect();



    drawLighting(srcPixelArray, absolutePaintSize.width(), absolutePaintSize.height());
}

#if ENABLE(OPENCL)
bool FELighting::platformApplyOpenCL()
{
    asImageBuffer();
    platformApplySoftware();
    ImageBuffer* sourceImage = asImageBuffer();
    RefPtr<Uint8ClampedArray> sourceImageData = sourceImage->getUnmultipliedImageData(IntRect(IntPoint(),sourceImage->internalSize()));
    createOpenCLImageResult(sourceImageData->data());
    return true;
}
#endif

} // namespace WebCore

#endif // ENABLE(FILTERS)

Lines 46-51 struct FELightingPaintingDataForNeon; a/Source/WebCore/platform/graphics/filters/FELighting.h_sec1
- a/Source/WebCore/platform/graphics/filters/FELighting.h +3 lines
46	class FELighting : public FilterEffect {	46	class FELighting : public FilterEffect {
47	public:	47	public:
48	virtual void platformApplySoftware();	48	virtual void platformApplySoftware();
		49	#if ENABLE(OPENCL)
		50	virtual bool platformApplyOpenCL();
		51	#endif
49	#if USE(SKIA)	52	#if USE(SKIA)
50	virtual bool platformApplySkia();	53	virtual bool platformApplySkia();
51	#endif	54	#endif

Lines 34-39 public: a/Source/WebCore/platform/graphics/filters/FEMerge.h_sec1
- a/Source/WebCore/platform/graphics/filters/FEMerge.h +3 lines
34	static PassRefPtr<FEMerge> create(Filter*);	34	static PassRefPtr<FEMerge> create(Filter*);
35		35
36	virtual void platformApplySoftware();	36	virtual void platformApplySoftware();
		37	#if ENABLE(OPENCL)
		38	virtual bool platformApplyOpenCL();
		39	#endif
37	virtual void dump();	40	virtual void dump();
38		41
39	virtual TextStream& externalRepresentation(TextStream&, int indention) const;	42	virtual TextStream& externalRepresentation(TextStream&, int indention) const;

Lines 47-52 public: a/Source/WebCore/platform/graphics/filters/FEMorphology.h_sec1
- a/Source/WebCore/platform/graphics/filters/FEMorphology.h +3 lines
47	bool setRadiusY(float);	47	bool setRadiusY(float);
48		48
49	virtual void platformApplySoftware();	49	virtual void platformApplySoftware();
		50	#if ENABLE(OPENCL)
		51	virtual bool platformApplyOpenCL();
		52	#endif
50	#if USE(SKIA)	53	#if USE(SKIA)
51	virtual bool platformApplySkia();	54	virtual bool platformApplySkia();
52	#endif	55	#endif

Lines 39-44 public: a/Source/WebCore/platform/graphics/filters/FEOffset.h_sec1
- a/Source/WebCore/platform/graphics/filters/FEOffset.h +3 lines
39	void setDy(float);	39	void setDy(float);
40		40
41	virtual void platformApplySoftware();	41	virtual void platformApplySoftware();
		42	#if ENABLE(OPENCL)
		43	virtual bool platformApplyOpenCL();
		44	#endif
42	virtual void dump();	45	virtual void dump();
43		46
44	virtual void determineAbsolutePaintRect();	47	virtual void determineAbsolutePaintRect();

Lines 33-38 public: a/Source/WebCore/platform/graphics/filters/FETile.h_sec1
- a/Source/WebCore/platform/graphics/filters/FETile.h +3 lines
33	static PassRefPtr<FETile> create(Filter* filter);	33	static PassRefPtr<FETile> create(Filter* filter);
34		34
35	virtual void platformApplySoftware();	35	virtual void platformApplySoftware();
		36	#if ENABLE(OPENCL)
		37	virtual bool platformApplyOpenCL();
		38	#endif
36	virtual void dump();	39	virtual void dump();
37		40
38	virtual void determineAbsolutePaintRect() { setAbsolutePaintRect(enclosingIntRect(maxEffectRect())); }	41	virtual void determineAbsolutePaintRect() { setAbsolutePaintRect(enclosingIntRect(maxEffectRect())); }

Lines 61-66 public: a/Source/WebCore/platform/graphics/filters/FETurbulence.h_sec1
- a/Source/WebCore/platform/graphics/filters/FETurbulence.h +3 lines
61	static void fillRegionWorker(void*);	61	static void fillRegionWorker(void*);
62		62
63	virtual void platformApplySoftware();	63	virtual void platformApplySoftware();
		64	#if ENABLE(OPENCL)
		65	virtual bool platformApplyOpenCL();
		66	#endif
64	virtual void dump();	67	virtual void dump();
65		68
66	virtual void determineAbsolutePaintRect() { setAbsolutePaintRect(enclosingIntRect(maxEffectRect())); }	69	virtual void determineAbsolutePaintRect() { setAbsolutePaintRect(enclosingIntRect(maxEffectRect())); }

Lines 36-42 a/Source/WebCore/platform/graphics/filters/FilterEffect.cpp_sec1
- a/Source/WebCore/platform/graphics/filters/FilterEffect.cpp +87 lines
36	namespace WebCore {	36	namespace WebCore {
37		37
38	FilterEffect::FilterEffect(Filter* filter)	38	FilterEffect::FilterEffect(Filter* filter)
		39	#if ENABLE(OPENCL)
		40	: m_openclImageResult(0)
		41	, m_alphaImage(false)
		42	#else
39	: m_alphaImage(false)	43	: m_alphaImage(false)
		44	#endif
40	, m_filter(filter)	45	, m_filter(filter)
41	, m_hasX(false)	46	, m_hasX(false)
42	, m_hasY(false)	47	, m_hasY(false)
Lines 51-56 FilterEffect::FilterEffect(Filter* filter) a/Source/WebCore/platform/graphics/filters/FilterEffect.cpp_sec2
51		56
52	FilterEffect::~FilterEffect()	57	FilterEffect::~FilterEffect()
53	{	58	{
		59	#if ENABLE(OPENCL)
		60	if (m_openclImageResult)
		61	clReleaseMemObject(m_openclImageResult);
		62	#endif
54	}	63	}
55		64
56	inline bool isFilterSizeValid(IntRect rect)	65	inline bool isFilterSizeValid(IntRect rect)
Lines 108-114 void FilterEffect::apply() a/Source/WebCore/platform/graphics/filters/FilterEffect.cpp_sec3
108	return;	117	return;
109		118
110	// Convert input results to the current effect's color space.	119	// Convert input results to the current effect's color space.
		120	#if ENABLE(OPENCL)
		121	in->transformResultColorSpaceOpenCL(m_colorSpace);
		122	#else
111	in->transformResultColorSpace(m_colorSpace);	123	in->transformResultColorSpace(m_colorSpace);
		124	#endif
112	}	125	}
113		126
114	determineAbsolutePaintRect();	127	determineAbsolutePaintRect();
Lines 123-128 void FilterEffect::apply() a/Source/WebCore/platform/graphics/filters/FilterEffect.cpp_sec4
123	}	136	}
124		137
125	// Add platform specific apply functions here and return earlier.	138	// Add platform specific apply functions here and return earlier.
		139	#if ENABLE(OPENCL)
		140	if (platformApplyOpenCL())
		141	return;
		142	#endif
126	#if USE(SKIA)	143	#if USE(SKIA)
127	if (platformApplySkia())	144	if (platformApplySkia())
128	return;	145	return;
Lines 187-192 void FilterEffect::clearResult() a/Source/WebCore/platform/graphics/filters/FilterEffect.cpp_sec5
187	m_unmultipliedImageResult.clear();	204	m_unmultipliedImageResult.clear();
188	if (m_premultipliedImageResult)	205	if (m_premultipliedImageResult)
189	m_premultipliedImageResult.clear();	206	m_premultipliedImageResult.clear();
		207	#if ENABLE(OPENCL)
		208	if (m_openclImageResult){
		209	clReleaseMemObject(m_openclImageResult);
		210	m_openclImageResult = 0;
		211	}
		212	#endif
190	}	213	}
191		214
192	ImageBuffer* FilterEffect::asImageBuffer()	215	ImageBuffer* FilterEffect::asImageBuffer()
Lines 195-200 ImageBuffer* FilterEffect::asImageBuffer() a/Source/WebCore/platform/graphics/filters/FilterEffect.cpp_sec6
195	return 0;	218	return 0;
196	if (m_imageBufferResult)	219	if (m_imageBufferResult)
197	return m_imageBufferResult.get();	220	return m_imageBufferResult.get();
		221	#if ENABLE(OPENCL)
		222	if (m_openclImageResult)
		223	return getOpenCLImageResult();
		224	#endif
198	m_imageBufferResult = ImageBuffer::create(m_absolutePaintRect.size(), 1, m_resultColorSpace, m_filter->renderingMode());	225	m_imageBufferResult = ImageBuffer::create(m_absolutePaintRect.size(), 1, m_resultColorSpace, m_filter->renderingMode());
199	IntRect destinationRect(IntPoint(), m_absolutePaintRect.size());	226	IntRect destinationRect(IntPoint(), m_absolutePaintRect.size());
200	if (m_premultipliedImageResult)	227	if (m_premultipliedImageResult)
Lines 204-209 ImageBuffer* FilterEffect::asImageBuffer() a/Source/WebCore/platform/graphics/filters/FilterEffect.cpp_sec7
204	return m_imageBufferResult.get();	231	return m_imageBufferResult.get();
205	}	232	}
206		233
		234	#if ENABLE(OPENCL)
		235	ImageBuffer* FilterEffect::getOpenCLImageResult()
		236	{
		237	FilterContextOpenCL* context = FilterContextOpenCL::context();
		238	ASSERT(context);
		239
		240	size_t origin[3] = { 0, 0, 0 };
		241	size_t region[3] = { m_absolutePaintRect.width(), m_absolutePaintRect.height(), 1};
		242
		243	RefPtr<Uint8ClampedArray> destinationPixelArray = Uint8ClampedArray::create(m_absolutePaintRect.width() * m_absolutePaintRect.height() * 4);
		244
		245	clEnqueueReadImage(context->commandQueue(), m_openclImageResult, CL_TRUE, origin, region, 0, 0, destinationPixelArray->data(), 0, 0, 0);
		246
		247	m_imageBufferResult = ImageBuffer::create(m_absolutePaintRect.size());
		248	IntRect destinationRect(IntPoint(), m_absolutePaintRect.size());
		249	m_imageBufferResult->putByteArray(Unmultiplied, destinationPixelArray.get(), destinationRect.size(), destinationRect, IntPoint());
		250
		251	return m_imageBufferResult.get();
		252	}
		253	#endif
		254
207	PassRefPtr<Uint8ClampedArray> FilterEffect::asUnmultipliedImage(const IntRect& rect)	255	PassRefPtr<Uint8ClampedArray> FilterEffect::asUnmultipliedImage(const IntRect& rect)
208	{	256	{
209	ASSERT(isFilterSizeValid(rect));	257	ASSERT(isFilterSizeValid(rect));
Lines 363-368 Uint8ClampedArray* FilterEffect::createPremultipliedImageResult() a/Source/WebCore/platform/graphics/filters/FilterEffect.cpp_sec8
363	return m_premultipliedImageResult.get();	411	return m_premultipliedImageResult.get();
364	}	412	}
365		413
		414	#if ENABLE(OPENCL)
		415	cl_mem FilterEffect::createOpenCLImageResult(unsigned char* source)
		416	{
		417	ASSERT(!hasResult());
		418	cl_image_format clImageFormat;
		419	clImageFormat.image_channel_order = CL_RGBA;
		420	clImageFormat.image_channel_data_type = CL_UNORM_INT8;
		421
		422	FilterContextOpenCL* context = FilterContextOpenCL::context();
		423	ASSERT(context);
		424	cl_int er;
		425	m_openclImageResult = clCreateImage2D(context->deviceContext(), source ? (CL_MEM_READ_WRITE \| CL_MEM_COPY_HOST_PTR) : CL_MEM_READ_WRITE,
		426	&clImageFormat, m_absolutePaintRect.width(), m_absolutePaintRect.height(), 0, source, &er);
		427	if (er != CL_SUCCESS)
		428	context->OpenCLPrintError(er);
		429	context->OpenCLDebugImage(m_openclImageResult);
		430	return m_openclImageResult;
		431	}
		432	#endif
		433
366	void FilterEffect::transformResultColorSpace(ColorSpace dstColorSpace)	434	void FilterEffect::transformResultColorSpace(ColorSpace dstColorSpace)
367	{	435	{
368	#if USE(CG)	436	#if USE(CG)
Lines 389-394 void FilterEffect::transformResultColorSpace(ColorSpace dstColorSpace) a/Source/WebCore/platform/graphics/filters/FilterEffect.cpp_sec9
389	#endif	457	#endif
390	}	458	}
391		459
		460	#if ENABLE(OPENCL)
		461	void FilterEffect::transformResultColorSpaceOpenCL(ColorSpace dstColorSpace)
		462	{
		463	if (!hasResult() \|\| dstColorSpace == m_resultColorSpace)
		464	return;
		465
		466	if (!m_openclImageResult) {
		467	asOpenCLImage();
		468	ASSERT(m_openclImageResult);
		469	}
		470
		471	FilterContextOpenCL* context = FilterContextOpenCL::context();
		472
		473	m_openclImageResult = context->OpenCLTransformColorSpace(m_openclImageResult, m_resultColorSpace, dstColorSpace, absolutePaintRect());
		474
		475	m_resultColorSpace = dstColorSpace;
		476	}
		477	#endif
		478
392	TextStream& FilterEffect::externalRepresentation(TextStream& ts, int) const	479	TextStream& FilterEffect::externalRepresentation(TextStream& ts, int) const
393	{	480	{
394	// FIXME: We should dump the subRegions of the filter primitives here later. This isn't	481	// FIXME: We should dump the subRegions of the filter primitives here later. This isn't

Lines 33-38 a/Source/WebCore/platform/graphics/filters/FilterEffect.h_sec1
- a/Source/WebCore/platform/graphics/filters/FilterEffect.h -1 / +25 lines
33	#include <wtf/Uint8ClampedArray.h>	33	#include <wtf/Uint8ClampedArray.h>
34	#include <wtf/Vector.h>	34	#include <wtf/Vector.h>
35		35
		36	#if ENABLE(OPENCL)
		37	#include "OpenCL/OpenCLContext.h"
		38	#endif
		39
36	static const float kMaxFilterSize = 5000.0f;	40	static const float kMaxFilterSize = 5000.0f;
37		41
38	namespace WebCore {	42	namespace WebCore {
Lines 62-67 public: a/Source/WebCore/platform/graphics/filters/FilterEffect.h_sec2
62	void copyUnmultipliedImage(Uint8ClampedArray* destination, const IntRect&);	66	void copyUnmultipliedImage(Uint8ClampedArray* destination, const IntRect&);
63	void copyPremultipliedImage(Uint8ClampedArray* destination, const IntRect&);	67	void copyPremultipliedImage(Uint8ClampedArray* destination, const IntRect&);
64		68
		69	#if ENABLE(OPENCL)
		70	cl_mem asOpenCLImage() { return m_openclImageResult; }
		71	ImageBuffer* getOpenCLImageResult();
		72	#endif
		73
65	FilterEffectVector& inputEffects() { return m_inputEffects; }	74	FilterEffectVector& inputEffects() { return m_inputEffects; }
66	FilterEffect* inputEffect(unsigned) const;	75	FilterEffect* inputEffect(unsigned) const;
67	unsigned numberOfEffectInputs() const { return m_inputEffects.size(); }	76	unsigned numberOfEffectInputs() const { return m_inputEffects.size(); }
Lines 69-75 public: a/Source/WebCore/platform/graphics/filters/FilterEffect.h_sec3
69	inline bool hasResult() const	78	inline bool hasResult() const
70	{	79	{
71	// This function needs platform specific checks, if the memory managment is not done by FilterEffect.	80	// This function needs platform specific checks, if the memory managment is not done by FilterEffect.
		81	#if ENABLE(OPENCL)
		82	return m_imageBufferResult \|\| m_unmultipliedImageResult \|\| m_premultipliedImageResult \|\| m_openclImageResult;
		83	#else
72	return m_imageBufferResult \|\| m_unmultipliedImageResult \|\| m_premultipliedImageResult;	84	return m_imageBufferResult \|\| m_unmultipliedImageResult \|\| m_premultipliedImageResult;
		85	#endif
73	}	86	}
74		87
75	IntRect drawingRegionOfInputImage(const IntRect&) const;	88	IntRect drawingRegionOfInputImage(const IntRect&) const;
Lines 93-98 public: a/Source/WebCore/platform/graphics/filters/FilterEffect.h_sec4
93	virtual void correctFilterResultIfNeeded() { }	106	virtual void correctFilterResultIfNeeded() { }
94		107
95	virtual void platformApplySoftware() = 0;	108	virtual void platformApplySoftware() = 0;
		109	#if ENABLE(OPENCL)
		110	virtual bool platformApplyOpenCL() = 0;
		111	#endif
96	#if USE(SKIA)	112	#if USE(SKIA)
97	virtual bool platformApplySkia() { return false; }	113	virtual bool platformApplySkia() { return false; }
98	#endif	114	#endif
Lines 133-138 public: a/Source/WebCore/platform/graphics/filters/FilterEffect.h_sec5
133	ColorSpace colorSpace() const { return m_colorSpace; }	149	ColorSpace colorSpace() const { return m_colorSpace; }
134	void setColorSpace(ColorSpace colorSpace) { m_colorSpace = colorSpace; }	150	void setColorSpace(ColorSpace colorSpace) { m_colorSpace = colorSpace; }
135	void transformResultColorSpace(ColorSpace);	151	void transformResultColorSpace(ColorSpace);
		152	#if ENABLE(OPENCL)
		153	void transformResultColorSpaceOpenCL(ColorSpace);
		154	#endif
136		155
137	protected:	156	protected:
138	FilterEffect(Filter*);	157	FilterEffect(Filter*);
Lines 140-145 protected: a/Source/WebCore/platform/graphics/filters/FilterEffect.h_sec6
140	ImageBuffer* createImageBufferResult();	159	ImageBuffer* createImageBufferResult();
141	Uint8ClampedArray* createUnmultipliedImageResult();	160	Uint8ClampedArray* createUnmultipliedImageResult();
142	Uint8ClampedArray* createPremultipliedImageResult();	161	Uint8ClampedArray* createPremultipliedImageResult();
		162	#if ENABLE(OPENCL)
		163	cl_mem createOpenCLImageResult(unsigned char* = 0);
		164	#endif
143		165
144	// Return true if the filter will only operate correctly on valid RGBA values, with	166	// Return true if the filter will only operate correctly on valid RGBA values, with
145	// alpha in [0,255] and each color component in [0, alpha].	167	// alpha in [0,255] and each color component in [0, alpha].
Lines 153-159 private: a/Source/WebCore/platform/graphics/filters/FilterEffect.h_sec7
153	RefPtr<Uint8ClampedArray> m_unmultipliedImageResult;	175	RefPtr<Uint8ClampedArray> m_unmultipliedImageResult;
154	RefPtr<Uint8ClampedArray> m_premultipliedImageResult;	176	RefPtr<Uint8ClampedArray> m_premultipliedImageResult;
155	FilterEffectVector m_inputEffects;	177	FilterEffectVector m_inputEffects;
156		178	#if ENABLE(OPENCL)
		179	cl_mem m_openclImageResult;
		180	#endif
157	bool m_alphaImage;	181	bool m_alphaImage;
158		182
159	IntRect m_absolutePaintRect;	183	IntRect m_absolutePaintRect;



#include "config.h"

#if ENABLE(OPENCL)
#include "OpenCLContext.h"
#include <iostream>
#include <time.h>
#include "stdio.h"

#define PROGRAM_STR(Src) #Src
#define PROGRAM(Src) PROGRAM_STR(Src)

namespace WebCore {

FilterContextOpenCL* FilterContextOpenCL::context()
{
     static bool wasInitialized = false;
     static FilterContextOpenCL* context = 0;

     if (wasInitialized)
         return context;

#ifdef NDEBUG
     printf("debug\n");
#endif
     wasInitialized = true;
     context = new FilterContextOpenCL();

     // Initializing the context.
     cl_int errNum;
     cl_device_id *devices;
     cl_platform_id firstPlatformId;
     size_t deviceBufferSize = -1;

     errNum = clGetPlatformIDs(1, &firstPlatformId, NULL);
     cl_context_properties contextProperties[] = { CL_CONTEXT_PLATFORM, (cl_context_properties)firstPlatformId, 0};
     context->m_deviceContext = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_GPU, NULL, NULL, &errNum);
     if (errNum != CL_SUCCESS) {
         context->m_deviceContext = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_CPU, NULL, NULL, &errNum);
         if (errNum != CL_SUCCESS)
            return context;
     }

     errNum = clGetContextInfo(context->m_deviceContext, CL_CONTEXT_DEVICES, 0, NULL, &deviceBufferSize);
     if (errNum != CL_SUCCESS)
         return context;

     if (deviceBufferSize <= 0)
         return context;

     devices = new cl_device_id[deviceBufferSize / sizeof(cl_device_id)];
     errNum = clGetContextInfo(context->m_deviceContext, CL_CONTEXT_DEVICES, deviceBufferSize, devices, NULL);
     if (errNum != CL_SUCCESS)
         return context;

     context->m_commandQueue = clCreateCommandQueue(context->m_deviceContext, devices[0], 0, NULL);
     if (context->m_commandQueue == NULL)
         return context;

     context->m_device = devices[0];
     delete [] devices;

     cl_bool imageSupport = CL_FALSE;
     clGetDeviceInfo(context->m_device, CL_DEVICE_IMAGE_SUPPORT, sizeof(cl_bool), &imageSupport, NULL);
     if (imageSupport != CL_TRUE)
         return context;

     return context;
}

cl_program FilterContextOpenCL::compileProgram(const char* source)
{
    cl_int errNum;
    cl_program program;

    FilterContextOpenCL* context = FilterContextOpenCL::context();

    program = clCreateProgramWithSource(context->m_deviceContext, 1, (const char**) &source, NULL, &errNum);
    if (errNum != CL_SUCCESS)
        OpenCLPrintError(errNum);
    if (errNum != CL_SUCCESS)
        return 0;

    errNum = clBuildProgram(program, 0, 0, 0, 0, 0);
    if (errNum != CL_SUCCESS)
        OpenCLPrintError(errNum);
    if (errNum != CL_SUCCESS)
        return 0;

    return program;
}

static const char* OpenCLTransformColorSpaceKernelProgram =
PROGRAM_STR(
    const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;

    __kernel void OpenCLTransformColorSpace(__read_only image2d_t source,
                                            __write_only image2d_t destination,
                                            __constant float *clLookUpTable)
    {
        int2 sourceCoord = (int2) (get_global_id(0), get_global_id(1));
        float4 pixel = read_imagef(source, sampler, sourceCoord);

        pixel = (float4)(clLookUpTable[(int)(round(pixel.x * 255))],
                         clLookUpTable[(int)(round(pixel.y * 255))],
                         clLookUpTable[(int)(round(pixel.z * 255))],
                         pixel.w);

        write_imagef(destination, sourceCoord, pixel);
    }
);

cl_mem FilterContextOpenCL::OpenCLTransformColorSpace(cl_mem source, ColorSpace resultColorSpace, ColorSpace dstColorSpace, IntRect sourceSize)
{
    DEFINE_STATIC_LOCAL(cl_mem, deviceRgbLUT, ());
    DEFINE_STATIC_LOCAL(cl_mem, linearRgbLUT, ());

    if ((resultColorSpace != ColorSpaceLinearRGB && resultColorSpace != ColorSpaceDeviceRGB)
        || (dstColorSpace != ColorSpaceLinearRGB && dstColorSpace != ColorSpaceDeviceRGB))
        return source;


    FilterContextOpenCL* context = FilterContextOpenCL::context();

    cl_image_format clImageFormat;
    clImageFormat.image_channel_order = CL_RGBA;
    clImageFormat.image_channel_data_type = CL_UNORM_INT8;

    cl_mem destination = clCreateImage2D(context->deviceContext(), CL_MEM_READ_WRITE,
        &clImageFormat, sourceSize.width(), sourceSize.height(), 0, 0, 0);

    if(!m_OpenCLTransformColorSpaceProgram){
        m_OpenCLTransformColorSpaceProgram = compileProgram(OpenCLTransformColorSpaceKernelProgram);
        m_OpenCLTransformColorSpace = kernelByName(m_OpenCLTransformColorSpaceProgram, "OpenCLTransformColorSpace");
    }

    RunKernel kernel(context, m_OpenCLTransformColorSpace, sourceSize.width(), sourceSize.height());
    kernel.addArgument(source);
    kernel.addArgument(destination);

    if (dstColorSpace == ColorSpaceLinearRGB) {
        if (!linearRgbLUT) {
            Vector<float> lookUpTable;
            for (unsigned i = 0; i < 256; i++) {
                float color = i  / 255.0f;
                color = (color <= 0.04045f ? color / 12.92f : pow((color + 0.055f) / 1.055f, 2.4f));
                color = std::max(0.0f, color);
                color = std::min(1.0f, color);
                lookUpTable.append((round(color * 255)) / 255);
            }
            linearRgbLUT = context->uploadBuffer(context, lookUpTable.data(), sizeof(float) * 256);
        }
        kernel.addArgument(linearRgbLUT);
    } else if (dstColorSpace == ColorSpaceDeviceRGB) {
        if (!deviceRgbLUT) {
            Vector<float> lookUpTable;
            for (unsigned i = 0; i < 256; i++) {
                float color = i / 255.0f;
                color = (powf(color, 1.0f / 2.4f) * 1.055f) - 0.055f;
                color = std::max(0.0f, color);
                color = std::min(1.0f, color);
                lookUpTable.append((round(color * 255)) / 255);
            }
            deviceRgbLUT = context->uploadBuffer(context, lookUpTable.data(), sizeof(float) * 256);
        }
        kernel.addArgument(deviceRgbLUT);
    }

    kernel.run();

    return 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);
            }
);

void FilterContextOpenCL::Fill(cl_mem image, IntSize imageSize, Color color)
{

    if (!m_Fill) {
        m_FillProgram = compileProgram(FillKernelProgram);
        ASSERT(m_FillProgram);
        m_Fill = kernelByName(m_FillProgram, "Fill");
        ASSERT(m_Fill);
    }

    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_mem FilterContextOpenCL::uploadBuffer(WebCore::FilterContextOpenCL* context, void* buffer, int size)
{
    cl_int err;
    cl_mem result = clCreateBuffer(context->deviceContext(), CL_MEM_READ_ONLY, size, 0, &err);
    if (err != CL_SUCCESS)
        OpenCLPrintError(err);

    err = clEnqueueWriteBuffer(context->commandQueue(), result, CL_TRUE, 0, size, buffer, 0, 0, 0);
    if (err != CL_SUCCESS)
        OpenCLPrintError(err);

    return result;
}

cl_mem FilterContextOpenCL::createOpenCLImage(FloatSize paintSize)
{
    FilterContextOpenCL* context = FilterContextOpenCL::context();

    cl_image_format clImageFormat;
    clImageFormat.image_channel_order = CL_RGBA;
    clImageFormat.image_channel_data_type = CL_UNORM_INT8;

    cl_int er;
    cl_mem image = clCreateImage2D(context->deviceContext(), CL_MEM_READ_WRITE ,
                           &clImageFormat, paintSize.width(), paintSize.height(), 0, 0, &er);
    if (er != CL_SUCCESS)
        context->OpenCLPrintError(er);
    context->OpenCLDebugImage(image);
    return image;
}

void FilterContextOpenCL::OpenCLPrintError(cl_int error)
{
    switch (error) {
    case CL_SUCCESS:
        fprintf(stderr, "CL_SUCCESS\n");
        break;
    case CL_DEVICE_NOT_FOUND:
        fprintf(stderr, "CL_DEVICE_NOT_FOUND\n");
        break;
    case CL_DEVICE_NOT_AVAILABLE:
        fprintf(stderr, "CL_DEVICE_NOT_AVAILABLE\n");
        break;
    case CL_COMPILER_NOT_AVAILABLE:
        fprintf(stderr, "CL_COMPILER_NOT_AVAILABLE\n");
        break;
    case CL_MEM_OBJECT_ALLOCATION_FAILURE:
        fprintf(stderr, "CL_MEM_OBJECT_ALLOCATION_FAILURE\n");
        break;
    case CL_OUT_OF_RESOURCES:
        fprintf(stderr, "CL_OUT_OF_RESOURCES\n");
        break;
    case CL_OUT_OF_HOST_MEMORY:
        fprintf(stderr, "CL_OUT_OF_HOST_MEMORY\n");
        break;
    case CL_PROFILING_INFO_NOT_AVAILABLE:
        fprintf(stderr, "CL_PROFILING_INFO_NOT_AVAILABLE\n");
        break;
    case CL_MEM_COPY_OVERLAP:
        fprintf(stderr, "CL_MEM_COPY_OVERLAP\n");
        break;
    case CL_IMAGE_FORMAT_MISMATCH:
        fprintf(stderr, "CL_IMAGE_FORMAT_MISMATCH\n");
        break;
    case CL_IMAGE_FORMAT_NOT_SUPPORTED:
        fprintf(stderr, "CL_IMAGE_FORMAT_NOT_SUPPORTED\n");
        break;
    case CL_BUILD_PROGRAM_FAILURE:
        fprintf(stderr, "CL_BUILD_PROGRAM_FAILURE\n");
        break;
    case CL_MAP_FAILURE:
        fprintf(stderr, "CL_MAP_FAILURE\n");
        break;
    case CL_INVALID_VALUE:
        fprintf(stderr, "CL_INVALID_VALUE\n");
        break;
    case CL_INVALID_DEVICE_TYPE:
        fprintf(stderr, "CL_INVALID_DEVICE_TYPE\n");
        break;
    case CL_INVALID_PLATFORM:
        fprintf(stderr, "CL_INVALID_PLATFORM\n");
        break;
    case CL_INVALID_DEVICE:
        fprintf(stderr, "CL_INVALID_DEVICE\n");
        break;
    case CL_INVALID_CONTEXT:
        fprintf(stderr, "CL_INVALID_CONTEXT\n");
        break;
    case CL_INVALID_QUEUE_PROPERTIES:
        fprintf(stderr, "CL_INVALID_QUEUE_PROPERTIES\n");
        break;
    case CL_INVALID_COMMAND_QUEUE:
        fprintf(stderr, "CL_INVALID_COMMAND_QUEUE\n");
        break;
    case CL_INVALID_HOST_PTR:
        fprintf(stderr, "CL_INVALID_HOST_PTR\n");
        break;
    case CL_INVALID_MEM_OBJECT:
        fprintf(stderr, "CL_INVALID_MEM_OBJECT\n");
        break;
    case CL_INVALID_IMAGE_FORMAT_DESCRIPTOR:
        fprintf(stderr, "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR\n");
        break;
    case CL_INVALID_IMAGE_SIZE:
        fprintf(stderr, "CL_INVALID_IMAGE_SIZE\n");
        break;
    case CL_INVALID_SAMPLER:
        fprintf(stderr, "CL_INVALID_SAMPLER\n");
        break;
    case CL_INVALID_BINARY:
        fprintf(stderr, "CL_INVALID_BINARY\n");
        break;
    case CL_INVALID_BUILD_OPTIONS:
        fprintf(stderr, "CL_INVALID_BUILD_OPTIONS\n");
        break;
    case CL_INVALID_PROGRAM:
        fprintf(stderr, "CL_INVALID_PROGRAM\n");
        break;
    case CL_INVALID_PROGRAM_EXECUTABLE:
        fprintf(stderr, "CL_INVALID_PROGRAM_EXECUTABLE\n");
        break;
    case CL_INVALID_KERNEL_NAME:
        fprintf(stderr, "CL_INVALID_KERNEL_NAME\n");
        break;
    case CL_INVALID_KERNEL_DEFINITION:
        fprintf(stderr, "CL_INVALID_KERNEL_DEFINITION\n");
        break;
    case CL_INVALID_KERNEL:
        fprintf(stderr, "CL_INVALID_KERNEL\n");
        break;
    case CL_INVALID_ARG_INDEX:
        fprintf(stderr, "CL_INVALID_ARG_INDEX\n");
        break;
    case CL_INVALID_ARG_VALUE:
        fprintf(stderr, "CL_INVALID_ARG_VALUE\n");
        break;
    case CL_INVALID_ARG_SIZE:
        fprintf(stderr, "CL_INVALID_ARG_SIZE\n");
        break;
    case CL_INVALID_KERNEL_ARGS:
        fprintf(stderr, "CL_INVALID_KERNEL_ARGS\n");
        break;
    case CL_INVALID_WORK_DIMENSION:
        fprintf(stderr, "CL_INVALID_WORK_DIMENSION\n");
        break;
    case CL_INVALID_WORK_GROUP_SIZE:
        fprintf(stderr, "CL_INVALID_WORK_GROUP_SIZE\n");
        break;
    case CL_INVALID_WORK_ITEM_SIZE:
        fprintf(stderr, "CL_INVALID_WORK_ITEM_SIZE\n");
        break;
    case CL_INVALID_GLOBAL_OFFSET:
        fprintf(stderr, "CL_INVALID_GLOBAL_OFFSET\n");
        break;
    case CL_INVALID_EVENT_WAIT_LIST:
        fprintf(stderr, "CL_INVALID_EVENT_WAIT_LIST\n");
        break;
    case CL_INVALID_EVENT:
        fprintf(stderr, "CL_INVALID_EVENT\n");
        break;
    case CL_INVALID_OPERATION:
        fprintf(stderr, "CL_INVALID_OPERATION\n");
        break;
    case CL_INVALID_GL_OBJECT:
        fprintf(stderr, "CL_INVALID_GL_OBJECT\n");
        break;
    case CL_INVALID_BUFFER_SIZE:
        fprintf(stderr, "CL_INVALID_BUFFER_SIZE\n");
        break;
    case CL_INVALID_MIP_LEVEL:
        fprintf(stderr, "CL_INVALID_MIP_LEVEL\n");
        break;
    case CL_INVALID_GLOBAL_WORK_SIZE:
        fprintf(stderr, "CL_INVALID_GLOBAL_WORK_SIZE\n");
        break;
    default:
        fprintf(stderr, "Unknown error code : %d\n", error);
        break;
    }
}

void FilterContextOpenCL::OpenCLDebugImage(cl_mem image)
{
    size_t width;
    size_t height;
    cl_int err;
    size_t channelSize;

    err = clGetImageInfo(image, CL_IMAGE_WIDTH, sizeof(width), &width, NULL);
    if (err != CL_SUCCESS)
        OpenCLPrintError(err);

    err = clGetImageInfo(image, CL_IMAGE_HEIGHT, sizeof(height), &height, NULL);
    if (err != CL_SUCCESS)
        OpenCLPrintError(err);

    err = clGetImageInfo(image, CL_IMAGE_ELEMENT_SIZE, sizeof(channelSize), &channelSize, NULL);
    if (err != CL_SUCCESS)
        OpenCLPrintError(err);
    if (err != CL_SUCCESS)
        printf("OpenCLDebugImage: width: %ld  height: %ld channelSize: %ld \n", width, height, channelSize);
}




} // namespace WebCore

#endif



#ifndef OpenCLContext_h
#define OpenCLContext_h

#include "config.h"
#if ENABLE(OPENCL)

#include "CL/cl.h"
#include "Color.h"
#include "ColorSpace.h"
#include "IntSize.h"
#include "IntRect.h"
#include "FloatPoint.h"
#include "FloatSize.h"
#include "stdio.h"

namespace WebCore {

class FilterContextOpenCL
{
public:
    // Returns NULL if initialization failed.
    static FilterContextOpenCL* context();

    cl_context deviceContext() { return m_deviceContext; }
    cl_command_queue commandQueue() { return m_commandQueue; }

    cl_mem createOpenCLImage(FloatSize);

    static cl_mem uploadBuffer(WebCore::FilterContextOpenCL* context, void* buffer, int size);
    cl_mem OpenCLTransformColorSpace(cl_mem, ColorSpace, ColorSpace, IntRect);
    void Fill(cl_mem, IntSize, Color);

    inline void compileFEComposite();
    inline void compileFEOffset();
    inline void compileFEBlend();
    inline void compileFEGaussianBlur();
    inline void compileFETurbulence();
    inline void compileFEColorMatrix();
    inline void compileFEConvolveMatrix();
    inline void compileFEComponentTransfer();
    inline void compileFEMerge();
    inline void compileFETile();
    inline void compileFEMorphology();

    inline void applyFEColorMatrix(cl_mem, IntSize, cl_mem, FloatPoint, cl_mem, int);
    inline void applyFEColorMatrixLuminanceToAlpha(cl_mem, IntSize, cl_mem, FloatPoint, int);
    inline void applyFEConvolveMatrix(cl_mem, IntSize, cl_mem, IntSize, cl_mem, int, float, float, IntPoint, bool);
    inline void applyFETurbulence(cl_mem, IntSize, cl_mem, cl_mem, cl_mem, cl_mem, cl_mem, cl_mem, int, int, int, int, float, float, bool, int, int);
    inline void applyFEComposite(cl_mem, IntSize, cl_mem, cl_mem, FloatPoint, FloatPoint, float, float, float, float, int);
    inline void applyFEComponentTransfer(cl_mem, IntSize, cl_mem, cl_mem, FloatPoint);
    inline void applyFEOffset(cl_mem, IntSize, cl_mem, FloatPoint);
    inline void applyFEBlend(cl_mem, IntSize, cl_mem, cl_mem, int,FloatPoint, FloatPoint);
    inline void applyFEGaussianBlur(cl_mem, IntSize, cl_mem);
    inline void applyFEMerge(cl_mem, IntSize, cl_mem, FloatRect);
    inline void applyFETiling(cl_mem, IntSize, cl_mem, FloatRect, FloatPoint);
    inline void applyFETileBuilder(cl_mem, FloatRect, cl_mem, FloatSize);
    inline void applyFEMorphology(cl_mem, IntSize, cl_mem, int, int, int, int, int);

    static void OpenCLPrintError(cl_int);
    static void OpenCLDebugImage(cl_mem);

private:

    class RunKernel {
    public:
        RunKernel(FilterContextOpenCL* context, cl_kernel kernel, size_t width, size_t height)
            : m_context(context)
            , m_kernel(kernel)
            , index(0)
        {
            m_globalSize[0] = width;
            m_globalSize[1] = height;
        }

        void addArgument(cl_mem handle)
        {
            cl_int err;
            err = clSetKernelArg(m_kernel, index++, sizeof(cl_mem), reinterpret_cast<void*>(&handle));
            if (err != CL_SUCCESS){
                m_context->OpenCLPrintError(err);
            }
        }

        void addArgument(cl_int value)
        {
            cl_int err;
            err = clSetKernelArg(m_kernel, index++, sizeof(cl_int), reinterpret_cast<void*>(&value));
            if (err != CL_SUCCESS){
                m_context->OpenCLPrintError(err);
            }
        }

        void addArgument(cl_float value)
        {
            cl_int err;
            err = clSetKernelArg(m_kernel, index++, sizeof(cl_float), reinterpret_cast<void*>(&value));
            if (err != CL_SUCCESS){
                m_context->OpenCLPrintError(err);
            }
        }

        void addArgument(cl_sampler handle)
        {
            cl_int err;
            err = clSetKernelArg(m_kernel, index++, sizeof(cl_sampler), reinterpret_cast<void*>(&handle));
            if (err != CL_SUCCESS){
                m_context->OpenCLPrintError(err);
            }
        }

        void run()
        {
            clFinish(m_context->m_commandQueue);
            clEnqueueNDRangeKernel(m_context->m_commandQueue, m_kernel, 2, 0, m_globalSize, 0, 0, 0, 0);
        }

        FilterContextOpenCL* m_context;
        cl_kernel m_kernel;
        size_t m_globalSize[2];
        int index;
    };

    static cl_program compileProgram(const char*);
    static inline cl_kernel kernelByName(cl_program program, const char* name) { return clCreateKernel(program, name, NULL); }

    cl_command_queue m_commandQueue;
    cl_device_id m_device;
    cl_context m_deviceContext;

    cl_program m_OpenCLTransformColorSpaceProgram;
    cl_kernel m_OpenCLTransformColorSpace;

    cl_program m_FEColorMatrixProgram;
    cl_kernel m_FEMatrix;
    cl_kernel m_FESaturate;
    cl_kernel m_FEHuerotate;
    cl_kernel m_FELuminance;

    cl_program m_FEConvolveMatrixProgram;
    cl_kernel m_FEConvolveMatrix;
    cl_kernel m_FEConvolveMatrixAlpha;

    cl_program m_FETurbulenceProgram;
    cl_kernel m_FETurbulence;

    cl_program m_feOffsetProgram;
    cl_kernel m_feOffset;

    cl_program m_feBlendProgram;
    cl_kernel m_feNormalBlend;
    cl_kernel m_feMultiplyBlend;
    cl_kernel m_feScreenBlend;
    cl_kernel m_feDarkenBlend;
    cl_kernel m_feLightenBlend;
    cl_kernel m_feUnknowBlend;

    cl_program m_feGaussianBlurProgram;
    cl_kernel m_feGaussianBlur;

    cl_program m_FEFloodProgram;
    cl_kernel m_FEFlood;

    cl_program m_feCompositeProgram;
    cl_kernel m_feArithmetic;
    cl_kernel m_feAtop;
    cl_kernel m_feIn;
    cl_kernel m_feOut;
    cl_kernel m_feOver;
    cl_kernel m_feXor;

    cl_program m_FEMergeProgram;
    cl_kernel m_FEMerge;

    cl_program m_feTileProgram;
    cl_kernel m_feTile;
    cl_kernel m_TileBuilder;

    cl_program m_feMorphologyProgram;
    cl_kernel m_feMorphologyDilate;
    cl_kernel m_feMorphologyErode;

    cl_program m_FillProgram;
    cl_kernel m_Fill;

    cl_program m_feComponentTransferProgram;
    cl_kernel m_feComponentTransfer;
};

} // namespace WebCore

#endif // ENABLE(OPENCL)

#endif



/*
 * Copyright (C) 2012 Tamas Czene <tczene@inf.u-szeged.hu>
 *
 * This library is free software; you can redistribute it and/or
 * modify it under the terms of the GNU Library General Public
 * License as published by the Free Software Foundation; either
 * version 2 of the License, or (at your option) any later version.
 *
 * This library is distributed in the hope that it will be useful,
 * but WITHOUT ANY WARRANTY; without even the implied warranty of
 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
 * Library General Public License for more details.
 *
 * You should have received a copy of the GNU Library General Public License
 * along with this library; see the file COPYING.LIB.  If not, write to
 * the Free Software Foundation, Inc., 51 Franklin Street, Fifth Floor,
 * Boston, MA 02110-1301, USA.
 */

#include "config.h"
#if ENABLE(FILTERS) && ENABLE(OPENCL)
#include "FEBlend.h"
#include "OpenCLContext.h"

#include <stdio.h>

#define PROGRAM_STR(Src) #Src
#define PROGRAM(Src) PROGRAM_STR(Src)

namespace WebCore {

static const char* FEBlendKernelProgram =
PROGRAM_STR(
            const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;

            __kernel void NormalBlend(__write_only image2d_t destination,
                                      __read_only image2d_t sourceA,
                                      __read_only image2d_t sourceB,
                                      float xA,
                                      float yA,
                                      float xB,
                                      float yB)
                {
                    int2 sourceCoordA = (int2) (get_global_id(0) + xA, get_global_id(1) + yA);
                    int2 sourceCoordB = (int2) (get_global_id(0) + xB, get_global_id(1) + yB);
                    float4 sourcePixelA = read_imagef(sourceA, sampler, sourceCoordA);
                    float4 sourcePixelB = read_imagef(sourceB, sampler, sourceCoordB);
                    float4 destinationPixel = (float4)((1.0 - sourcePixelA.w) * sourcePixelB.x + sourcePixelA.x,
                                                       (1.0 - sourcePixelA.w) * sourcePixelB.y + sourcePixelA.y,
                                                       (1.0 - sourcePixelA.w) * sourcePixelB.z + sourcePixelA.z,
                                                        1.0 - (1.0 - sourcePixelA.w) * (1.0 - sourcePixelB.w));
                    write_imagef(destination, (int2)(get_global_id(0), get_global_id(1)), destinationPixel);
                }

            __kernel void MultiplyBlend(__write_only image2d_t destination,
                                        __read_only image2d_t sourceA,
                                        __read_only image2d_t sourceB,
                                        float xA,
                                        float yA,
                                        float xB,
                                        float yB)
                {
                    int2 sourceCoordA = (int2) (get_global_id(0) + xA, get_global_id(1) + yA);
                    int2 sourceCoordB = (int2) (get_global_id(0) + xB, get_global_id(1) + yB);
                    float4 sourcePixelA = read_imagef(sourceA, sampler, sourceCoordA);
                    float4 sourcePixelB = read_imagef(sourceB, sampler, sourceCoordB);
                    float4 destinationPixel = (float4)((1.0 - sourcePixelA.w) * sourcePixelB.x + (1.0 - sourcePixelB.w) * sourcePixelA.x + sourcePixelA.x * sourcePixelB.x,
                                                       (1.0 - sourcePixelA.w) * sourcePixelB.y + (1.0 - sourcePixelB.w) * sourcePixelA.y + sourcePixelA.y * sourcePixelB.y,
                                                       (1.0 - sourcePixelA.w) * sourcePixelB.z + (1.0 - sourcePixelB.w) * sourcePixelA.z + sourcePixelA.z * sourcePixelB.z,
                                                        1.0 - (1.0 - sourcePixelA.w) * (1.0 - sourcePixelB.w));
                    write_imagef(destination, (int2)(get_global_id(0), get_global_id(1)), destinationPixel);
                }

            __kernel void ScreenBlend(__write_only image2d_t destination,
                                      __read_only image2d_t sourceA,
                                      __read_only image2d_t sourceB,
                                      float xA,
                                      float yA,
                                      float xB,
                                      float yB)
                {
                    int2 sourceCoordA = (int2) (get_global_id(0) + xA, get_global_id(1) + yA);
                    int2 sourceCoordB = (int2) (get_global_id(0) + xB, get_global_id(1) + yB);
                    float4 sourcePixelA = read_imagef(sourceA, sampler, sourceCoordA);
                    float4 sourcePixelB = read_imagef(sourceB, sampler, sourceCoordB);
                    float4 destinationPixel = (float4)(sourcePixelB.x + sourcePixelA.x - sourcePixelA.x * sourcePixelB.x,
                                                       sourcePixelB.y + sourcePixelA.y - sourcePixelA.y * sourcePixelB.y,
                                                       sourcePixelB.z + sourcePixelA.z - sourcePixelA.z * sourcePixelB.z,
                                                       1.0 - (1.0 - sourcePixelA.w) * (1.0 - sourcePixelB.w));
                    write_imagef(destination, (int2)(get_global_id(0), get_global_id(1)), destinationPixel);
                }

            __kernel void DarkenBlend(__write_only image2d_t destination,
                                      __read_only image2d_t sourceA,
                                      __read_only image2d_t sourceB,
                                      float xA,
                                      float yA,
                                      float xB,
                                      float yB)
                {
                    int2 sourceCoordA = (int2) (get_global_id(0) + xA, get_global_id(1) + yA);
                    int2 sourceCoordB = (int2) (get_global_id(0) + xB, get_global_id(1) + yB);
                    float4 sourcePixelA = read_imagef(sourceA, sampler, sourceCoordA);
                    float4 sourcePixelB = read_imagef(sourceB, sampler, sourceCoordB);
                    float4 destinationPixel = (float4)(min((1.0 - sourcePixelA.w) * sourcePixelB.x + sourcePixelA.x, (1.0 - sourcePixelB.w) * sourcePixelA.x + sourcePixelB.x),
                                                       min((1.0 - sourcePixelA.w) * sourcePixelB.y + sourcePixelA.y, (1.0 - sourcePixelB.w) * sourcePixelA.y + sourcePixelB.y),
                                                       min((1.0 - sourcePixelA.w) * sourcePixelB.z + sourcePixelA.z, (1.0 - sourcePixelB.w) * sourcePixelA.z + sourcePixelB.z),
                                                            1.0 - (1.0 - sourcePixelA.w) * (1.0 - sourcePixelB.w));
                    write_imagef(destination, (int2)(get_global_id(0), get_global_id(1)), destinationPixel);
                }

            __kernel void LightenBlend(__write_only image2d_t destination,
                                       __read_only image2d_t sourceA,
                                       __read_only image2d_t sourceB,
                                       float xA,
                                       float yA,
                                       float xB,
                                       float yB)
                {
                    int2 sourceCoordA = (int2) (get_global_id(0) + xA, get_global_id(1) + yA);
                    int2 sourceCoordB = (int2) (get_global_id(0) + xB, get_global_id(1) + yB);
                    float4 sourcePixelA = read_imagef(sourceA, sampler, sourceCoordA);
                    float4 sourcePixelB = read_imagef(sourceB, sampler, sourceCoordB);
                    float4 destinationPixel = (float4)(max((1.0 - sourcePixelA.w) * sourcePixelB.x + sourcePixelA.x, (1.0 - sourcePixelB.w) * sourcePixelA.x + sourcePixelB.x),
                                                       max((1.0 - sourcePixelA.w) * sourcePixelB.y + sourcePixelA.y, (1.0 - sourcePixelB.w) * sourcePixelA.y + sourcePixelB.y),
                                                       max((1.0 - sourcePixelA.w) * sourcePixelB.z + sourcePixelA.z, (1.0 - sourcePixelB.w) * sourcePixelA.z + sourcePixelB.z),
                                                            1.0 - (1.0 - sourcePixelA.w) * (1.0 - sourcePixelB.w));
                    write_imagef(destination, (int2)(get_global_id(0), get_global_id(1)), destinationPixel);
                }

            __kernel void UnknowBlend(__write_only image2d_t destination,
                                      __read_only image2d_t sourceA,
                                      __read_only image2d_t sourceB,
                                      float xA,
                                      float yA,
                                      float xB,
                                      float yB)
                {
                    int2 sourceCoord = (int2) (get_global_id(0), get_global_id(1));
                    float4 destinationPixel = (float4)(0.0,0.0,0.0,0.0);
                    write_imagef(destination, (int2)(get_global_id(0), get_global_id(1)), destinationPixel);
                }
);

inline void FilterContextOpenCL::compileFEBlend()
{
    if (!m_feNormalBlend) {
        m_feBlendProgram = compileProgram(FEBlendKernelProgram);
        ASSERT(m_feBlendProgram);

        m_feNormalBlend = kernelByName(m_feBlendProgram, "NormalBlend");
        ASSERT(m_feNormalBlend);
        m_feMultiplyBlend = kernelByName(m_feBlendProgram, "MultiplyBlend");
        ASSERT(m_feNormalBlend);
        m_feScreenBlend = kernelByName(m_feBlendProgram, "ScreenBlend");
        ASSERT(m_feNormalBlend);
        m_feDarkenBlend = kernelByName(m_feBlendProgram, "DarkenBlend");
        ASSERT(m_feNormalBlend);
        m_feLightenBlend = kernelByName(m_feBlendProgram, "LightenBlend");
        ASSERT(m_feNormalBlend);
        m_feUnknowBlend = kernelByName(m_feBlendProgram, "UnknowBlend");
        ASSERT(m_feNormalBlend);
    }
}

inline void FilterContextOpenCL::applyFEBlend(cl_mem destination, IntSize destinationSize, cl_mem sourceA, cl_mem sourceB, int mode, FloatPoint relativeSourceLocationA, FloatPoint relativeSourceLocationB)
{
    cl_kernel blendMode;

    switch (mode) {
    case FEBLEND_MODE_NORMAL:
        blendMode = m_feNormalBlend;
        break;
    case FEBLEND_MODE_MULTIPLY:
        blendMode = m_feMultiplyBlend;
        break;
    case FEBLEND_MODE_SCREEN:
        blendMode = m_feScreenBlend;
        break;
    case FEBLEND_MODE_DARKEN:
        blendMode = m_feDarkenBlend;
        break;
    case FEBLEND_MODE_LIGHTEN:
        blendMode = m_feLightenBlend;
        break;
    case FEBLEND_MODE_UNKNOWN:
    default:
        blendMode = m_feUnknowBlend;
        break;
    }

    float xA = relativeSourceLocationA.x();
    float yA = relativeSourceLocationA.y();

    float xB = relativeSourceLocationB.x();
    float yB = relativeSourceLocationB.y();

    RunKernel kernel(this, blendMode, destinationSize.width(), destinationSize.height());
    kernel.addArgument(destination);
    kernel.addArgument(sourceA);
    kernel.addArgument(sourceB);
    kernel.addArgument(xA);
    kernel.addArgument(yA);
    kernel.addArgument(xB);
    kernel.addArgument(yB);
    kernel.run();
}

bool FEBlend::platformApplyOpenCL()
{
    FilterContextOpenCL* context = FilterContextOpenCL::context();

    if (!context)
        return false;

    context->compileFEBlend();

    FilterEffect* in = inputEffect(0);
    FilterEffect* in2 = inputEffect(1);

    FloatPoint relativeSourceLocationA(in->absolutePaintRect().location());
    relativeSourceLocationA.setX(absolutePaintRect().x() - relativeSourceLocationA.x());
    relativeSourceLocationA.setY(absolutePaintRect().y() - relativeSourceLocationA.y());

    FloatPoint relativeSourceLocationB(in2->absolutePaintRect().location());
    relativeSourceLocationB.setX(absolutePaintRect().x() - relativeSourceLocationB.x());
    relativeSourceLocationB.setY(absolutePaintRect().y() - relativeSourceLocationB.y());

    cl_mem sourceA = in->asOpenCLImage();
    cl_mem sourceB = in2->asOpenCLImage();
    cl_mem destination = createOpenCLImageResult();

    context->applyFEBlend(destination, absolutePaintRect().size(), sourceA, sourceB, m_mode, relativeSourceLocationA, relativeSourceLocationB);

    return true;
}

} // namespace WebCore

#endif // ENABLE(FILTERS) && ENABLE(OPENCL)



/*
 * Copyright (C) 2012 Tamas Czene <tczene@inf.u-szeged.hu>
 *
 * This library is free software; you can redistribute it and/or
 * modify it under the terms of the GNU Library General Public
 * License as published by the Free Software Foundation; either
 * version 2 of the License, or (at your option) any later version.
 *
 * This library is distributed in the hope that it will be useful,
 * but WITHOUT ANY WARRANTY; without even the implied warranty of
 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
 * Library General Public License for more details.
 *
 * You should have received a copy of the GNU Library General Public License
 * along with this library; see the file COPYING.LIB.  If not, write to
 * the Free Software Foundation, Inc., 51 Franklin Street, Fifth Floor,
 * Boston, MA 02110-1301, USA.
 */

#include "config.h"
#if ENABLE(FILTERS) && ENABLE(OPENCL)
#include "FEColorMatrix.h"
#include "OpenCLContext.h"
#include "stdio.h"

#define PROGRAM_STR(Src) #Src
#define PROGRAM(Src) PROGRAM_STR(Src)

namespace WebCore {

static const char* FEColorMatrixKernelProgram =
PROGRAM_STR(
    const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
    __constant float pi = 3.14159265358979323846;

    __kernel void Matrix(__write_only image2d_t destination,
                         __read_only image2d_t source,
                         float x,
                         float y,
                         __constant float *values)
    {
        int2 sourceCoord = (int2) (get_global_id(0) + x, get_global_id(1) + y);
        float4 sourcePixel = read_imagef(source, sampler, sourceCoord);

        float4 destinationPixel = (float4)(values[0]  * sourcePixel.x + values[1]  * sourcePixel.y + values[2]  * sourcePixel.z + values[3]  * sourcePixel.w + values[4] * 255,
                                             values[5]  * sourcePixel.x + values[6]  * sourcePixel.y + values[7]  * sourcePixel.z + values[8]  * sourcePixel.w + values[9] * 255,
                                             values[10] * sourcePixel.x + values[11] * sourcePixel.y + values[12] * sourcePixel.z + values[13] * sourcePixel.w + values[14] * 255,
                                             values[15] * sourcePixel.x + values[16] * sourcePixel.y + values[17] * sourcePixel.z + values[18] * sourcePixel.w + values[19] * 255);

        write_imagef(destination, (int2)(get_global_id(0), get_global_id(1)), destinationPixel);
    }


    __kernel void Saturate(__write_only image2d_t destination,
                         __read_only image2d_t source,
                         float x,
                         float y,
                         __constant float *values)
    {
        int2 sourceCoord = (int2) (get_global_id(0) + x, get_global_id(1) + y);
        float4 sourcePixel = read_imagef(source, sampler, sourceCoord);

        float4 destinationPixel = (float4)((0.213 + 0.787 * values[0]) * sourcePixel.x + (0.715 - 0.715 * values[0]) * sourcePixel.y + (0.072 - 0.072 * values[0]) * sourcePixel.z,
                                           (0.213 - 0.213 * values[0]) * sourcePixel.x + (0.715 + 0.285 * values[0]) * sourcePixel.y + (0.072 - 0.072 * values[0]) * sourcePixel.z,
                                           (0.213 - 0.213 * values[0]) * sourcePixel.x + (0.715 - 0.715 * values[0]) * sourcePixel.y + (0.072 + 0.928 * values[0]) * sourcePixel.z,
                                           sourcePixel.w);

        write_imagef(destination, (int2)(get_global_id(0), get_global_id(1)), destinationPixel);
    }

    __kernel void Huerotate(__write_only image2d_t destination,
                         __read_only image2d_t source,
                         float x,
                         float y,
                         __constant float *values)
    {
        int2 sourceCoord = (int2) (get_global_id(0) + x, get_global_id(1) + y);
        float4 sourcePixel = read_imagef(source, sampler, sourceCoord);

        float cosHue = cos(values[0] * pi / 180);
        float sinHue = sin(values[0] * pi / 180);

        float4 destinationPixel = (float4)(sourcePixel.x * (0.213 + cosHue * 0.787 - sinHue * 0.213) + sourcePixel.y * (0.715 - cosHue * 0.715 - sinHue * 0.715) + sourcePixel.z  * (0.072 - cosHue * 0.072 + sinHue * 0.928),
                                           sourcePixel.x * (0.213 - cosHue * 0.213 + sinHue * 0.143) + sourcePixel.y * (0.715 + cosHue * 0.285 + sinHue * 0.140) + sourcePixel.z  * (0.072 - cosHue * 0.072 - sinHue * 0.283),
                                           sourcePixel.x * (0.213 - cosHue * 0.213 - sinHue * 0.787) + sourcePixel.y * (0.715 - cosHue * 0.715 + sinHue * 0.715) + sourcePixel.z  * (0.072 + cosHue * 0.928 + sinHue * 0.072),
                                           sourcePixel.w);

        write_imagef(destination, (int2)(get_global_id(0), get_global_id(1)), destinationPixel);
    }

    __kernel void Luminance(__write_only image2d_t destination,
                         __read_only image2d_t source,
                         float x,
                         float y)
    {
        int2 sourceCoord = (int2) (get_global_id(0) + x, get_global_id(1) + y);
        float4 sourcePixel = read_imagef(source, sampler, sourceCoord);

        float4 destinationPixel = (float4)(0, 0, 0, 0.2125 * sourcePixel.x + 0.7154 * sourcePixel.y + 0.0721 * sourcePixel.z);

        write_imagef(destination, (int2)(get_global_id(0), get_global_id(1)), sourcePixel);
    }


);

inline void FilterContextOpenCL::compileFEColorMatrix()
{
    if (!m_FEColorMatrixProgram) {
        m_FEColorMatrixProgram = compileProgram(FEColorMatrixKernelProgram);
        ASSERT(m_FEColorMatrixProgram);

        m_FEMatrix = kernelByName(m_FEColorMatrixProgram, "Matrix");
        ASSERT(m_FEMatrix);
        m_FESaturate = kernelByName(m_FEColorMatrixProgram, "Saturate");
        ASSERT(m_FESaturate);
        m_FEHuerotate = kernelByName(m_FEColorMatrixProgram, "Huerotate");
        ASSERT(m_FEHuerotate);
        m_FELuminance = kernelByName(m_FEColorMatrixProgram, "Luminance");
        ASSERT(m_FELuminance);
    }
}

inline void FilterContextOpenCL::applyFEColorMatrix(cl_mem destination, IntSize destinationSize, cl_mem source, FloatPoint relativeSourceLocation, cl_mem values, int type)
{
    float x = relativeSourceLocation.x();
    float y = relativeSourceLocation.y();

    cl_kernel m_FEColorMatrix;

    switch (type) {
    case FECOLORMATRIX_TYPE_MATRIX:
        m_FEColorMatrix = m_FEMatrix;
        break;
    case FECOLORMATRIX_TYPE_SATURATE:
        m_FEColorMatrix = m_FESaturate;
        break;
    case FECOLORMATRIX_TYPE_HUEROTATE:
        m_FEColorMatrix = m_FEHuerotate;
        break;
    case FECOLORMATRIX_TYPE_LUMINANCETOALPHA:
        m_FEColorMatrix = m_FELuminance;
        break;
    default:
        m_FEColorMatrix = 0;
    }

    RunKernel kernel(this, m_FEColorMatrix, destinationSize.width(), destinationSize.height());
    kernel.addArgument(destination);
    kernel.addArgument(source);
    kernel.addArgument(x);
    kernel.addArgument(y);
    kernel.addArgument(values);
    kernel.run();

}

inline void FilterContextOpenCL::applyFEColorMatrixLuminanceToAlpha(cl_mem destination, IntSize destinationSize, cl_mem source, FloatPoint relativeSourceLocation, int m_type)
{
    float x = relativeSourceLocation.x();
    float y = relativeSourceLocation.y();

    RunKernel kernel(this, m_FELuminance, destinationSize.width(), destinationSize.height());
    kernel.addArgument(destination);
    kernel.addArgument(source);
    kernel.addArgument(x);
    kernel.addArgument(y);
    kernel.run();
}

bool FEColorMatrix::platformApplyOpenCL()
{
    FilterContextOpenCL* context = FilterContextOpenCL::context();

    if (!context)
        return false;

    context->compileFEColorMatrix();

    FilterEffect* in = inputEffect(0);
    cl_mem source = in->asOpenCLImage();
    cl_mem destination = createOpenCLImageResult();
    cl_mem clValues = 0;

    if(m_type != FECOLORMATRIX_TYPE_LUMINANCETOALPHA)
        clValues = context->uploadBuffer(context, m_values.data(), sizeof(float) * 20);

    FloatPoint relativeSourceLocation(in->absolutePaintRect().location());
    relativeSourceLocation.setX(absolutePaintRect().x() - relativeSourceLocation.x());
    relativeSourceLocation.setY(absolutePaintRect().y() - relativeSourceLocation.y());

    if(m_type != FECOLORMATRIX_TYPE_LUMINANCETOALPHA)
        context->applyFEColorMatrix(destination, absolutePaintRect().size(), source, relativeSourceLocation, clValues, m_type);
    else
        context->applyFEColorMatrixLuminanceToAlpha(destination, absolutePaintRect().size(), source, relativeSourceLocation, m_type);

    clReleaseMemObject(clValues);
    return true;
}



} // namespace WebCore

#endif // ENABLE(FILTERS) && ENABLE(OPENCL)



/*
 * Copyright (C) 2012 Tamas Czene <tczene@inf.u-szeged.hu>
 *
 * This library is free software; you can redistribute it and/or
 * modify it under the terms of the GNU Library General Public
 * License as published by the Free Software Foundation; either
 * version 2 of the License, or (at your option) any later version.
 *
 * This library is distributed in the hope that it will be useful,
 * but WITHOUT ANY WARRANTY; without even the implied warranty of
 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
 * Library General Public License for more details.
 *
 * You should have received a copy of the GNU Library General Public License
 * along with this library; see the file COPYING.LIB.  If not, write to
 * the Free Software Foundation, Inc., 51 Franklin Street, Fifth Floor,
 * Boston, MA 02110-1301, USA.
 */

#include "config.h"
#if ENABLE(FILTERS) && ENABLE(OPENCL)
#include "FEComposite.h"
#include "OpenCLContext.h"

#include <stdio.h>

#define PROGRAM_STR(Src) #Src
#define PROGRAM(Src) PROGRAM_STR(Src)

namespace WebCore {

static const char* FECompositeKernelProgram =
PROGRAM_STR(
    const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;

    __kernel void Arithmetic(__write_only image2d_t destination,
                         __read_only image2d_t sourceA,
                         __read_only image2d_t sourceB,
                         float xA,
                         float yA,
                         float xB,
                         float yB,
                         float k1,
                         float k2,
                         float k3,
                         float k4)
    {
        int2 sourceCoordA = (int2) (get_global_id(0) + xA, get_global_id(1) + yA);
        int2 sourceCoordB = (int2) (get_global_id(0) + xB, get_global_id(1) + yB);
        float4 sourcePixelA = read_imagef(sourceA, sampler, sourceCoordA);
        float4 sourcePixelB = read_imagef(sourceB, sampler, sourceCoordB);
        sourcePixelA.xyz *= sourcePixelA.w;
        sourcePixelB.xyz *= sourcePixelB.w;

        float4 destinationPixel = (float4)(k1 * sourcePixelA.x * sourcePixelB.x + k2 * sourcePixelA.x + k3 * sourcePixelB.x + k4 ,
                                           k1 * sourcePixelA.y * sourcePixelB.y + k2 * sourcePixelA.y + k3 * sourcePixelB.y + k4 ,
                                           k1 * sourcePixelA.z * sourcePixelB.z + k2 * sourcePixelA.z + k3 * sourcePixelB.z + k4 ,
                                           k1 * sourcePixelA.w * sourcePixelB.w + k2 * sourcePixelA.w + k3 * sourcePixelB.w + k4);

        destinationPixel.w = clamp(destinationPixel.w, 0.0, 1.0);
        float rec = 1 / destinationPixel.w;
        destinationPixel.x = clamp(destinationPixel.x, 0.0, 1.0) * rec;
        destinationPixel.y = clamp(destinationPixel.y, 0.0, 1.0) * rec;
        destinationPixel.z = clamp(destinationPixel.z, 0.0, 1.0) * rec;

        write_imagef(destination, (int2) (get_global_id(0), get_global_id(1)), destinationPixel);
    }

    __kernel void Over(__write_only image2d_t destination,
                         __read_only image2d_t sourceA,
                         __read_only image2d_t sourceB,
                         float xA,
                         float yA,
                         float xB,
                         float yB)
    {
        int2 sourceCoordA = (int2) (get_global_id(0) + xA, get_global_id(1) + yA);
        int2 sourceCoordB = (int2) (get_global_id(0) + xB, get_global_id(1) + yB);
        float4 sourcePixelA = read_imagef(sourceA, sampler, sourceCoordA);
        float4 sourcePixelB = read_imagef(sourceB, sampler, sourceCoordB);
        sourcePixelA.xyz *= sourcePixelA.w;
        sourcePixelB.xyz *= sourcePixelB.w;

        float4 destinationPixel = (float4)(sourcePixelA.x + sourcePixelB.x * (1 - sourcePixelA.w),
                                           sourcePixelA.y + sourcePixelB.y * (1 - sourcePixelA.w),
                                           sourcePixelA.z + sourcePixelB.z * (1 - sourcePixelA.w),
                                           sourcePixelA.w + sourcePixelB.w * (1 - sourcePixelA.w));

        destinationPixel.w = clamp(destinationPixel.w, 0.0, 1.0);
        float rec = 1 / destinationPixel.w;
        destinationPixel.x = clamp(destinationPixel.x, 0.0, 1.0) * rec;
        destinationPixel.y = clamp(destinationPixel.y, 0.0, 1.0) * rec;
        destinationPixel.z = clamp(destinationPixel.z, 0.0, 1.0) * rec;

        write_imagef(destination, (int2) (get_global_id(0), get_global_id(1)), destinationPixel);
    }

    __kernel void In(__write_only image2d_t destination,
                         __read_only image2d_t sourceA,
                         __read_only image2d_t sourceB,
                         float xA,
                         float yA,
                         float xB,
                         float yB)
    {
        int2 sourceCoordA = (int2) (get_global_id(0) + xA, get_global_id(1) + yA);
        int2 sourceCoordB = (int2) (get_global_id(0) + xB, get_global_id(1) + yB);
        float4 sourcePixelA = read_imagef(sourceA, sampler, sourceCoordA);
        float4 sourcePixelB = read_imagef(sourceB, sampler, sourceCoordB);
        sourcePixelA.xyz *= sourcePixelA.w;
        sourcePixelB.xyz *= sourcePixelB.w;

        float4 destinationPixel = (float4)(sourcePixelA.x * sourcePixelB.w,
                                           sourcePixelA.y * sourcePixelB.w,
                                           sourcePixelA.z * sourcePixelB.w,
                                           sourcePixelA.w * sourcePixelB.w);

        destinationPixel.w = clamp(destinationPixel.w, 0.0, 1.0);
        float rec = 1 / destinationPixel.w;
        destinationPixel.x = clamp(destinationPixel.x, 0.0, 1.0) * rec;
        destinationPixel.y = clamp(destinationPixel.y, 0.0, 1.0) * rec;
        destinationPixel.z = clamp(destinationPixel.z, 0.0, 1.0) * rec;

        write_imagef(destination, (int2) (get_global_id(0), get_global_id(1)), destinationPixel);
    }

    __kernel void Out(__write_only image2d_t destination,
                         __read_only image2d_t sourceA,
                         __read_only image2d_t sourceB,
                         float xA,
                         float yA,
                         float xB,
                         float yB)
    {
        int2 sourceCoordA = (int2) (get_global_id(0) + xA, get_global_id(1) + yA);
        int2 sourceCoordB = (int2) (get_global_id(0) + xB, get_global_id(1) + yB);
        float4 sourcePixelA = read_imagef(sourceA, sampler, sourceCoordA);
        float4 sourcePixelB = read_imagef(sourceB, sampler, sourceCoordB);
        sourcePixelA.xyz *= sourcePixelA.w;
        sourcePixelB.xyz *= sourcePixelB.w;

        float4 destinationPixel = (float4)(sourcePixelA.x * (1 - sourcePixelB.w),
                                           sourcePixelA.y * (1 - sourcePixelB.w),
                                           sourcePixelA.z * (1 - sourcePixelB.w),
                                           sourcePixelA.w * (1 - sourcePixelB.w));

        destinationPixel.w = clamp(destinationPixel.w, 0.0, 1.0);
        float rec = 1 / destinationPixel.w;
        destinationPixel.x = clamp(destinationPixel.x, 0.0, 1.0) * rec;
        destinationPixel.y = clamp(destinationPixel.y, 0.0, 1.0) * rec;
        destinationPixel.z = clamp(destinationPixel.z, 0.0, 1.0) * rec;

        write_imagef(destination, (int2) (get_global_id(0), get_global_id(1)), destinationPixel);
    }

    __kernel void Atop(__write_only image2d_t destination,
                         __read_only image2d_t sourceA,
                         __read_only image2d_t sourceB,
                         float xA,
                         float yA,
                         float xB,
                         float yB)
    {
        int2 sourceCoordA = (int2) (get_global_id(0) + xA, get_global_id(1) + yA);
        int2 sourceCoordB = (int2) (get_global_id(0) + xB, get_global_id(1) + yB);
        float4 sourcePixelA = read_imagef(sourceA, sampler, sourceCoordA);
        float4 sourcePixelB = read_imagef(sourceB, sampler, sourceCoordB);
        sourcePixelA.xyz *= sourcePixelA.w;
        sourcePixelB.xyz *= sourcePixelB.w;

        float4 destinationPixel = (float4)(sourcePixelA.x * sourcePixelB.w + sourcePixelB.x * (1 - sourcePixelA.w),
                                           sourcePixelA.y * sourcePixelB.w + sourcePixelB.y * (1 - sourcePixelA.w),
                                           sourcePixelA.z * sourcePixelB.w + sourcePixelB.z * (1 - sourcePixelA.w),
                                           sourcePixelA.w * sourcePixelB.w + sourcePixelB.w * (1 - sourcePixelA.w));

        destinationPixel.w = clamp(destinationPixel.w, 0.0, 1.0);
        float rec = (destinationPixel.w == 0) ? 1.0 :  1.0 / destinationPixel.w;
        destinationPixel.x = clamp(destinationPixel.x, 0.0, 1.0) * rec;
        destinationPixel.y = clamp(destinationPixel.y, 0.0, 1.0) * rec;
        destinationPixel.z = clamp(destinationPixel.z, 0.0, 1.0) * rec;

        write_imagef(destination, (int2) (get_global_id(0), get_global_id(1)), destinationPixel);
    }

    __kernel void Xor(__write_only image2d_t destination,
                         __read_only image2d_t sourceA,
                         __read_only image2d_t sourceB,
                         float xA,
                         float yA,
                         float xB,
                         float yB)
    {
        int2 sourceCoordA = (int2) (get_global_id(0) + xA, get_global_id(1) + yA);
        int2 sourceCoordB = (int2) (get_global_id(0) + xB, get_global_id(1) + yB);
        float4 sourcePixelA = read_imagef(sourceA, sampler, sourceCoordA);
        float4 sourcePixelB = read_imagef(sourceB, sampler, sourceCoordB);
        sourcePixelA.xyz *= sourcePixelA.w;
        sourcePixelB.xyz *= sourcePixelB.w;

        float4 destinationPixel = (float4)(sourcePixelA.x * (1 - sourcePixelB.w) + sourcePixelB.x * (1 - sourcePixelA.w),
                                           sourcePixelA.y * (1 - sourcePixelB.w) + sourcePixelB.y * (1 - sourcePixelA.w),
                                           sourcePixelA.z * (1 - sourcePixelB.w) + sourcePixelB.z * (1 - sourcePixelA.w),
                                           sourcePixelA.w * (1 - sourcePixelB.w) + sourcePixelB.w * (1 - sourcePixelA.w));

        destinationPixel.w = clamp(destinationPixel.w, 0.0, 1.0);
        float rec = 1 / destinationPixel.w;
        destinationPixel.x = clamp(destinationPixel.x, 0.0, 1.0) * rec;
        destinationPixel.y = clamp(destinationPixel.y, 0.0, 1.0) * rec;
        destinationPixel.z = clamp(destinationPixel.z, 0.0, 1.0) * rec;

        write_imagef(destination, (int2) (get_global_id(0), get_global_id(1)), destinationPixel);
    }

);

inline void FilterContextOpenCL::compileFEComposite()
{
    if (!m_feCompositeProgram) {
        m_feCompositeProgram = compileProgram(FECompositeKernelProgram);
        ASSERT(m_feCompositeProgram);

        m_feArithmetic = kernelByName(m_feCompositeProgram, "Arithmetic");
        ASSERT(m_feArithmetic);
        m_feOver = kernelByName(m_feCompositeProgram, "Over");
        ASSERT(m_feOver);
        m_feIn = kernelByName(m_feCompositeProgram, "In");
        ASSERT(m_feIn);
        m_feOut = kernelByName(m_feCompositeProgram, "Out");
        ASSERT(m_feOut);
        m_feAtop = kernelByName(m_feCompositeProgram, "Atop");
        ASSERT(m_feAtop);
        m_feXor = kernelByName(m_feCompositeProgram, "Xor");
        ASSERT(m_feXor);
    }
}

inline void FilterContextOpenCL::applyFEComposite(cl_mem destination, IntSize destinationSize, cl_mem sourceA, cl_mem sourceB, FloatPoint relativeSourceLocationA, FloatPoint relativeSourceLocationB, float k1, float k2, float k3, float k4, int mode)
{
    cl_kernel compositeMode;

    switch (mode) {
    case FECOMPOSITE_OPERATOR_ARITHMETIC:
        compositeMode = m_feArithmetic;
        break;
    case FECOMPOSITE_OPERATOR_OVER:
        compositeMode = m_feOver;
        break;
    case FECOMPOSITE_OPERATOR_IN:
        compositeMode = m_feIn;
        break;
    case FECOMPOSITE_OPERATOR_OUT:
        compositeMode = m_feOut;
        break;
    case FECOMPOSITE_OPERATOR_ATOP:
        compositeMode = m_feAtop;
        break;
    case FECOMPOSITE_OPERATOR_XOR:
        compositeMode = m_feXor;
        break;
    default:
        compositeMode = m_feOver;
        break;
    }

    float xA = relativeSourceLocationA.x();
    float yA = relativeSourceLocationA.y();

    float xB = relativeSourceLocationB.x();
    float yB = relativeSourceLocationB.y();

    RunKernel kernel(this, compositeMode, destinationSize.width(), destinationSize.height());
    kernel.addArgument(destination);
    kernel.addArgument(sourceA);
    kernel.addArgument(sourceB);
    kernel.addArgument(xA);
    kernel.addArgument(yA);
    kernel.addArgument(xB);
    kernel.addArgument(yB);
    if (mode == FECOMPOSITE_OPERATOR_ARITHMETIC) {
        kernel.addArgument(k1);
        kernel.addArgument(k2);
        kernel.addArgument(k3);
        kernel.addArgument(k4);
    }
    kernel.run();

}

bool FEComposite::platformApplyOpenCL()
{
    FilterContextOpenCL* context = FilterContextOpenCL::context();

    if (!context)
        return false;

    context->compileFEComposite();

    FilterEffect* in = inputEffect(0);
    cl_mem sourceA = in->asOpenCLImage();

    FilterEffect* in2 = inputEffect(1);
    cl_mem sourceB = in2->asOpenCLImage();

    cl_mem destination = createOpenCLImageResult();

    FloatPoint relativeSourceLocationA(in->absolutePaintRect().location());
    relativeSourceLocationA.setX(absolutePaintRect().x() - relativeSourceLocationA.x());
    relativeSourceLocationA.setY(absolutePaintRect().y() - relativeSourceLocationA.y());

    FloatPoint relativeSourceLocationB(in2->absolutePaintRect().location());
    relativeSourceLocationB.setX(absolutePaintRect().x() - relativeSourceLocationB.x());
    relativeSourceLocationB.setY(absolutePaintRect().y() - relativeSourceLocationB.y());

    context->applyFEComposite(destination, absolutePaintRect().size(), sourceA, sourceB, relativeSourceLocationA, relativeSourceLocationB, m_k1, m_k2, m_k3, m_k4, m_type);

    return true;
}

} // namespace WebCore

#endif // ENABLE(FILTERS) && ENABLE(OPENCL)



/*
 * Copyright (C) 2012 Tamas Czene <tczene@inf.u-szeged.hu>
 *
 * This library is free software; you can redistribute it and/or
 * modify it under the terms of the GNU Library General Public
 * License as published by the Free Software Foundation; either
 * version 2 of the License, or (at your option) any later version.
 *
 * This library is distributed in the hope that it will be useful,
 * but WITHOUT ANY WARRANTY; without even the implied warranty of
 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
 * Library General Public License for more details.
 *
 * You should have received a copy of the GNU Library General Public License
 * along with this library; see the file COPYING.LIB.  If not, write to
 * the Free Software Foundation, Inc., 51 Franklin Street, Fifth Floor,
 * Boston, MA 02110-1301, USA.
 */

#include "config.h"
#if ENABLE(FILTERS) && ENABLE(OPENCL)
#include "FEConvolveMatrix.h"
#include "OpenCLContext.h"
#include "SVGFilter.h"
#include <fstream>
#include <stdio.h>

#define PROGRAM_STR(Src) #Src
#define PROGRAM(Src) PROGRAM_STR(Src)

namespace WebCore {

static const char* FEConvolveMatrixKernelProgram =
PROGRAM_STR(
__kernel void ConvolveMatrix(__write_only image2d_t destination, __read_only image2d_t source, __constant float *kernelMatrix,
                            int kernelSizeWidth, int kernelSizeHeight, int kernelOffsetX, int kernelOffsetY, float divisor,
                            float bias, sampler_t sampler)
{
    int XCoord = get_global_id(0) - kernelOffsetX;
    int YCoord = get_global_id(1) - kernelOffsetY;
    float4 totals = (float4)(0,0,0,0);
    int kernelIndex = kernelSizeWidth * kernelSizeHeight - 1;
    for(int y = 0; y < kernelSizeHeight; ++y) {
       for(int x = 0; x < kernelSizeWidth; ++x) {
           float4 sourcePixel = read_imagef(source, sampler, (int2)(XCoord + x, YCoord + y));
           sourcePixel.xzy *= sourcePixel.w;
           totals += sourcePixel * kernelMatrix[kernelIndex];
           --kernelIndex;
       }
    }
    totals = totals / divisor + bias;
    totals.w = clamp(totals.w, 0.0f, 1.0f);
    totals.xyz = clamp(totals.xyz, 0.0f, totals.w) / totals.w;
    write_imagef(destination, (int2)(get_global_id(0), get_global_id(1)), totals);
}

__kernel void ConvolveMatrixAlpha(__write_only image2d_t destination, __read_only image2d_t source, __constant float *kernelMatrix,
                                int kernelSizeWidth, int kernelSizeHeight, int kernelOffsetX, int kernelOffsetY, float divisor,
                                float bias, sampler_t sampler)
{
    int XCoord = get_global_id(0) - kernelOffsetX;
    int YCoord = get_global_id(1) - kernelOffsetY;
    float4 totals = (float4)(0,0,0,0);
    int kernelIndex = kernelSizeWidth * kernelSizeHeight - 1;
    for(int y = 0; y < kernelSizeHeight; ++y) {
       for(int x = 0; x < kernelSizeWidth; ++x) {
           float4 sourcePixel = read_imagef(source, sampler, (int2)(XCoord + x, YCoord + y));
           totals.xyz += sourcePixel.xyz * kernelMatrix[kernelIndex];
           --kernelIndex;
       }
    }
    float4 alpha = read_imagef(source, sampler, (int2)(get_global_id(0), get_global_id(1)));
    totals.xyz = clamp(totals.xyz / divisor + bias, 0.0f, 1.0f);
    totals.w = alpha.w;
    write_imagef(destination, (int2)(get_global_id(0), get_global_id(1)), totals);
}
);

inline void FilterContextOpenCL::compileFEConvolveMatrix()
{

    if (!m_FEConvolveMatrix) {
        m_FEConvolveMatrixProgram = compileProgram(FEConvolveMatrixKernelProgram);
        ASSERT(m_FEConvolveMatrixProgram);
        m_FEConvolveMatrix = kernelByName(m_FEConvolveMatrixProgram, "ConvolveMatrix");
        ASSERT(m_FEConvolveMatrix);
        m_FEConvolveMatrixAlpha = kernelByName(m_FEConvolveMatrixProgram, "ConvolveMatrixAlpha");
        ASSERT(m_FEConvolveMatrixAlpha);
    }
}

inline void FilterContextOpenCL::applyFEConvolveMatrix(cl_mem destination, IntSize destinationSize, cl_mem source, IntSize kernelSize, cl_mem kernelMatrix, int edgeMode, float divisor, float bias, IntPoint targetOffset, bool preserveAlpha)
{
    cl_sampler sampler;

    switch (edgeMode) {
    case EDGEMODE_DUPLICATE:
        sampler = clCreateSampler(deviceContext(), CL_FALSE, CL_ADDRESS_REPEAT, CL_FILTER_NEAREST, 0);
        break;
    case EDGEMODE_WRAP:
        sampler = clCreateSampler(deviceContext(), CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, 0);
        break;
    default:
        sampler = clCreateSampler(deviceContext(), CL_FALSE, CL_ADDRESS_NONE, CL_FILTER_NEAREST, 0);
    }

    int kernelOffsetX = targetOffset.x();
    int kernelOffsetY = targetOffset.y();

    int kernelSizeWidth = kernelSize.width();
    int kernelSizeHeight = kernelSize.height();

    RunKernel kernel(this, preserveAlpha ? m_FEConvolveMatrixAlpha : m_FEConvolveMatrix, destinationSize.width(), destinationSize.height());
    kernel.addArgument(destination);
    kernel.addArgument(source);
    kernel.addArgument(kernelMatrix);
    kernel.addArgument(kernelSizeWidth);
    kernel.addArgument(kernelSizeHeight);
    kernel.addArgument(kernelOffsetX);
    kernel.addArgument(kernelOffsetY);
    kernel.addArgument(divisor);
    kernel.addArgument(bias);
    kernel.addArgument(sampler);
    kernel.run();
}

bool FEConvolveMatrix::platformApplyOpenCL()
{
    FilterContextOpenCL* context = FilterContextOpenCL::context();
    if (!context)
        return false;

    context->compileFEConvolveMatrix();

    FilterEffect* in = inputEffect(0);

    cl_mem source = in->asOpenCLImage();
    cl_mem destination = createOpenCLImageResult();

    cl_mem kernelMatrix = context->uploadBuffer(context, m_kernelMatrix.data(), sizeof(float) * m_kernelSize.width() * m_kernelSize.height());

    context->applyFEConvolveMatrix(destination, absolutePaintRect().size(), source, m_kernelSize, kernelMatrix, m_edgeMode, m_divisor, m_bias, m_targetOffset, m_preserveAlpha);

    return true;
}

} // namespace WebCore

#endif // ENABLE(FILTERS) && ENABLE(OPENCL)



/*
 * Copyright (C) 2012 Tamas Czene <tczene@inf.u-szeged.hu>
 *
 * This library is free software; you can redistribute it and/or
 * modify it under the terms of the GNU Library General Public
 * License as published by the Free Software Foundation; either
 * version 2 of the License, or (at your option) any later version.
 *
 * This library is distributed in the hope that it will be useful,
 * but WITHOUT ANY WARRANTY; without even the implied warranty of
 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
 * Library General Public License for more details.
 *
 * You should have received a copy of the GNU Library General Public License
 * along with this library; see the file COPYING.LIB.  If not, write to
 * the Free Software Foundation, Inc., 51 Franklin Street, Fifth Floor,
 * Boston, MA 02110-1301, USA.
 */

#include "config.h"
#if ENABLE(FILTERS) && ENABLE(OPENCL)
#include "FEFlood.h"
#include "OpenCLContext.h"

#include <stdio.h>

#define PROGRAM_STR(Src) #Src
#define PROGRAM(Src) PROGRAM_STR(Src)

namespace WebCore {

bool FEFlood::platformApplyOpenCL()
{
    FilterContextOpenCL* context = FilterContextOpenCL::context();

    if (!context)
        return false;

    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)



/*
 * Copyright (C) 2012 Tamas Czene <tczene@inf.u-szeged.hu>
 *
 * This library is free software; you can redistribute it and/or
 * modify it under the terms of the GNU Library General Public
 * License as published by the Free Software Foundation; either
 * version 2 of the License, or (at your option) any later version.
 *
 * This library is distributed in the hope that it will be useful,
 * but WITHOUT ANY WARRANTY; without even the implied warranty of
 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
 * Library General Public License for more details.
 *
 * You should have received a copy of the GNU Library General Public License
 * along with this library; see the file COPYING.LIB.  If not, write to
 * the Free Software Foundation, Inc., 51 Franklin Street, Fifth Floor,
 * Boston, MA 02110-1301, USA.
 */

#include "config.h"
#if ENABLE(FILTERS) && ENABLE(OPENCL)

#include "OpenCLContext.h"
#include "SVGFEImage.h"
#include "SVGFilter.h"
#include "SharedBuffer.h"

#define PROGRAM_STR(Src) #Src
#define PROGRAM(Src) PROGRAM_STR(Src)

namespace WebCore {

bool FEImage::platformApplyOpenCL()
{
    platformApplySoftware();
    createOpenCLImageResult(asPremultipliedImage(absolutePaintRect())->data());
    return true;
}

} // namespace WebCore

#endif // ENABLE(FILTERS) && ENABLE(OPENCL)



/*
 * Copyright (C) 2012 Tamas Czene <tczene@inf.u-szeged.hu>
 *
 * This library is free software; you can redistribute it and/or
 * modify it under the terms of the GNU Library General Public
 * License as published by the Free Software Foundation; either
 * version 2 of the License, or (at your option) any later version.
 *
 * This library is distributed in the hope that it will be useful,
 * but WITHOUT ANY WARRANTY; without even the implied warranty of
 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
 * Library General Public License for more details.
 *
 * You should have received a copy of the GNU Library General Public License
 * along with this library; see the file COPYING.LIB.  If not, write to
 * the Free Software Foundation, Inc., 51 Franklin Street, Fifth Floor,
 * Boston, MA 02110-1301, USA.
 */

#include "config.h"
#if ENABLE(FILTERS) && ENABLE(OPENCL)
#include "FEMerge.h"
#include "OpenCLContext.h"

#define PROGRAM_STR(Src) #Src
#define PROGRAM(Src) PROGRAM_STR(Src)

namespace WebCore {

static const char* FEMergeKernelProgram =
PROGRAM_STR(
    __kernel void Merge(__write_only image2d_t destination,
                      __read_only image2d_t source,
                      float x,
                      float y,
                      float width,
                      float height,
                      sampler_t sampler)
    {
        int2 sourceCoord = (int2) (get_global_id(0) + x, get_global_id(1) + y);
        float4 sourcePixel = read_imagef(source, sampler, sourceCoord);
        if(sourceCoord.x >= 0 && sourceCoord.y >= 0 && sourceCoord.x < width && sourceCoord.y < height)
            write_imagef(destination, (int2) (get_global_id(0), get_global_id(1)), sourcePixel);
    }

);

inline void FilterContextOpenCL::compileFEMerge()
{
    if (!m_FEMergeProgram) {
        m_FEMergeProgram = compileProgram(FEMergeKernelProgram);
        ASSERT(m_FEMergeProgram);
        m_FEMerge = kernelByName(m_FEMergeProgram, "Merge");
        ASSERT(m_FEMerge);
    }
}

inline void FilterContextOpenCL::applyFEMerge(cl_mem destination, IntSize destinationSize, cl_mem source, FloatRect relativeSourceRect)
{
    float x = relativeSourceRect.x();
    float y = relativeSourceRect.y();
    float width = relativeSourceRect.width();
    float height = relativeSourceRect.height();

    cl_sampler sampler = clCreateSampler(deviceContext(), CL_FALSE, CL_ADDRESS_CLAMP,
                                         (round(x) == x && round(y) == y) ? CL_FILTER_NEAREST : CL_FILTER_LINEAR, 0);

    RunKernel kernel(this, m_FEMerge, destinationSize.width(), destinationSize.height());
    kernel.addArgument(destination);
    kernel.addArgument(source);
    kernel.addArgument(x);
    kernel.addArgument(y);
    kernel.addArgument(width);
    kernel.addArgument(height);
    kernel.addArgument(sampler);
    kernel.run();

}

bool FEMerge::platformApplyOpenCL()
{
    FilterContextOpenCL* context = FilterContextOpenCL::context();

    if (!context)
        return false;

    context->compileFEMerge();

    unsigned size = numberOfEffectInputs();
    ASSERT(size > 0);

    cl_mem destination = createOpenCLImageResult();
    cl_mem source;
    FilterEffect* in;

    for (unsigned i = 0; i < size; ++i){
        in = inputEffect(i);
        source = in->asOpenCLImage();
        FloatRect relativeSourceRect(in->absolutePaintRect());
        relativeSourceRect.setX(absolutePaintRect().x() - relativeSourceRect.x());
        relativeSourceRect.setY(absolutePaintRect().y() - relativeSourceRect.y());

        context->applyFEMerge(destination, absolutePaintRect().size(), source, relativeSourceRect);
    }

    return true;
}



} // namespace WebCore

#endif // ENABLE(FILTERS) && ENABLE(OPENCL)



/*
 * Copyright (C) 2012 Tamas Czene <tczene@inf.u-szeged.hu>
 *
 * This library is free software; you can redistribute it and/or
 * modify it under the terms of the GNU Library General Public
 * License as published by the Free Software Foundation; either
 * version 2 of the License, or (at your option) any later version.
 *
 * This library is distributed in the hope that it will be useful,
 * but WITHOUT ANY WARRANTY; without even the implied warranty of
 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
 * Library General Public License for more details.
 *
 * You should have received a copy of the GNU Library General Public License
 * along with this library; see the file COPYING.LIB.  If not, write to
 * the Free Software Foundation, Inc., 51 Franklin Street, Fifth Floor,
 * Boston, MA 02110-1301, USA.
 */

#include "config.h"
#if ENABLE(FILTERS) && ENABLE(OPENCL)
#include "FEMorphology.h"
#include "OpenCLContext.h"
#include "stdio.h"

#define PROGRAM_STR(Src) #Src
#define PROGRAM(Src) PROGRAM_STR(Src)

using std::min;

namespace WebCore {

static const char* FEMorphologyKernelProgram =
PROGRAM_STR(
    __kernel void MorphologyDilate(__write_only image2d_t destination,
                     __read_only image2d_t source,
                      int width,
                      int height,
                      int radiusX,
                      int radiusY,
                      sampler_t sampler)
    {
        int coordX = get_global_id(0);
        int coordY = get_global_id(1);

        int startX = max(0, coordX - radiusX);
        int endX = min(width, coordX + radiusX);

        int startY = max(0, coordY - radiusY);
        int endY = min(height, coordY + radiusY);

        float4 sourcePixel = (float4)(0,0,0,0);

        for(int y = startY; y <= endY; y++)
            for(int x = startX; x <= endX; x++){
                sourcePixel = max(sourcePixel , read_imagef(source, sampler, (int2)(x, y)));
            }

        write_imagef(destination, (int2)(get_global_id(0), get_global_id(1)), sourcePixel);
    }

    __kernel void MorphologyErode(__write_only image2d_t destination,
                     __read_only image2d_t source,
                      int width,
                      int height,
                      int radiusX,
                      int radiusY,
                      sampler_t sampler)
    {
        int coordX = get_global_id(0);
        int coordY = get_global_id(1);

        int startX = max(0, coordX - radiusX);
        int endX = min(width - 1, coordX + radiusX);

        int startY = max(0, coordY - radiusY);
        int endY = min(height - 1, coordY + radiusY);

        float4 sourcePixel = (float4)(256,256,256,256);

        for(int y = startY; y <= endY; y++)
            for(int x = startX; x <= endX; x++){
                sourcePixel = min(sourcePixel , read_imagef(source, sampler, (int2)(x, y)));
            }

        write_imagef(destination, (int2)(get_global_id(0), get_global_id(1)), sourcePixel);
    }

);

inline void FilterContextOpenCL::compileFEMorphology()
{
    if (!m_feMorphologyDilate && !m_feMorphologyErode) {
        m_feMorphologyProgram = compileProgram(FEMorphologyKernelProgram);
        ASSERT(m_feMorphologyProgram);
        m_feMorphologyDilate = kernelByName(m_feMorphologyProgram, "MorphologyDilate");
        ASSERT(m_feMorphologyDilate);
        m_feMorphologyErode = kernelByName(m_feMorphologyProgram, "MorphologyErode");
        ASSERT(m_feMorphologyErode);
    }
}

inline void FilterContextOpenCL::applyFEMorphology(cl_mem destination, IntSize destinationSize, cl_mem source, int width, int height, int radiusX, int radiusY, int type)
{

    cl_sampler sampler = clCreateSampler(deviceContext(), CL_FALSE, CL_ADDRESS_NONE, CL_FILTER_NEAREST, 0);

    RunKernel kernel(this, (type == 1) ? m_feMorphologyErode : m_feMorphologyDilate, destinationSize.width(), destinationSize.height());

    kernel.addArgument(destination);
    kernel.addArgument(source);
    kernel.addArgument(width);
    kernel.addArgument(height);
    kernel.addArgument(radiusX);
    kernel.addArgument(radiusY);
    kernel.addArgument(sampler);
    kernel.run();
}

bool FEMorphology::platformApplyOpenCL()
{
    FilterContextOpenCL* context = FilterContextOpenCL::context();

    if (!context)
        return false;

    context->compileFEMorphology();

    FilterEffect* in = inputEffect(0);
    cl_mem source = in->asOpenCLImage();
    cl_mem destination = createOpenCLImageResult();

    setIsAlphaImage(in->isAlphaImage());
    if (m_radiusX <= 0 || m_radiusY <= 0) {
        context->Fill(destination, absolutePaintRect().size(), Color(0,0,0,0));
        return true;
    }

    Filter* filter = this->filter();
    IntRect effectDrawingRect = requestedRegionOfInputImageData(in->absolutePaintRect());
    int width = effectDrawingRect.width();
    int height = effectDrawingRect.height();
    int radiusX = min(width - 1, static_cast<int>(floorf(filter->applyHorizontalScale(m_radiusX))));
    int radiusY = min(height - 1, static_cast<int>(floorf(filter->applyVerticalScale(m_radiusY))));

    context->applyFEMorphology(destination, absolutePaintRect().size(), source, width, height, radiusX, radiusY, m_type);

    return true;
}

} // namespace WebCore

#endif // ENABLE(FILTERS) && ENABLE(OPENCL)



/*
 * Copyright (C) 2012 Tamas Czene <tczene@inf.u-szeged.hu>
 *
 * This library is free software; you can redistribute it and/or
 * modify it under the terms of the GNU Library General Public
 * License as published by the Free Software Foundation; either
 * version 2 of the License, or (at your option) any later version.
 *
 * This library is distributed in the hope that it will be useful,
 * but WITHOUT ANY WARRANTY; without even the implied warranty of
 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
 * Library General Public License for more details.
 *
 * You should have received a copy of the GNU Library General Public License
 * along with this library; see the file COPYING.LIB.  If not, write to
 * the Free Software Foundation, Inc., 51 Franklin Street, Fifth Floor,
 * Boston, MA 02110-1301, USA.
 */

#include "config.h"
#if ENABLE(FILTERS) && ENABLE(OPENCL)
#include "FEOffset.h"
#include "OpenCLContext.h"
#include "stdio.h"

#define PROGRAM_STR(Src) #Src
#define PROGRAM(Src) PROGRAM_STR(Src)

namespace WebCore {

static const char* FEOffsetKernelProgram =
PROGRAM_STR(
    __kernel void Offset(__write_only image2d_t destination,
                     __read_only image2d_t source,
                      float x,
                      float y,
                      sampler_t sampler)
    {
        int2 sourceCoord = (int2) (get_global_id(0) + x , get_global_id(1) + y);
        float4 sourcePixel = read_imagef(source, sampler, sourceCoord);
        write_imagef(destination, (int2)(get_global_id(0), get_global_id(1)), sourcePixel);
    }

);

inline void FilterContextOpenCL::compileFEOffset()
{
    if (!m_feOffset) {
        m_feOffsetProgram = compileProgram(FEOffsetKernelProgram);
        ASSERT(m_feOffsetProgram);
        m_feOffset = kernelByName(m_feOffsetProgram, "Offset");
        ASSERT(m_feOffset);
    }
}

inline void FilterContextOpenCL::applyFEOffset(cl_mem destination, IntSize destinationSize, cl_mem source, FloatPoint relativeSourceLocation)
{

    float x = relativeSourceLocation.x();
    float y = relativeSourceLocation.y();

    cl_sampler sampler = clCreateSampler(deviceContext(), CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE,
                                         (round(x) == x && round(y) == y) ? CL_FILTER_NEAREST : CL_FILTER_LINEAR, 0);

    RunKernel kernel(this, m_feOffset, destinationSize.width(), destinationSize.height());
    kernel.addArgument(destination);
    kernel.addArgument(source);
    kernel.addArgument(x);
    kernel.addArgument(y);
    kernel.addArgument(sampler);
    kernel.run();
}

bool FEOffset::platformApplyOpenCL()
{
    FilterContextOpenCL* context = FilterContextOpenCL::context();

    if (!context)
        return false;

    context->compileFEOffset();

    FilterEffect* in = inputEffect(0);
    cl_mem source = in->asOpenCLImage();
    cl_mem destination = createOpenCLImageResult();

    FloatPoint relativeSourceLocation(in->absolutePaintRect().location());
    relativeSourceLocation.setX(absolutePaintRect().x() - (relativeSourceLocation.x() + filter()->applyHorizontalScale(m_dx)));
    relativeSourceLocation.setY(absolutePaintRect().y() - (relativeSourceLocation.y() + filter()->applyVerticalScale(m_dy)));
    context->applyFEOffset(destination, absolutePaintRect().size(), source,relativeSourceLocation);

    return true;
}

} // namespace WebCore

#endif // ENABLE(FILTERS) && ENABLE(OPENCL)



/*
 * Copyright (C) 2012 Tamas Czene <tczene@inf.u-szeged.hu>
 *
 * This library is free software; you can redistribute it and/or
 * modify it under the terms of the GNU Library General Public
 * License as published by the Free Software Foundation; either
 * version 2 of the License, or (at your option) any later version.
 *
 * This library is distributed in the hope that it will be useful,
 * but WITHOUT ANY WARRANTY; without even the implied warranty of
 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
 * Library General Public License for more details.
 *
 * You should have received a copy of the GNU Library General Public License
 * along with this library; see the file COPYING.LIB.  If not, write to
 * the Free Software Foundation, Inc., 51 Franklin Street, Fifth Floor,
 * Boston, MA 02110-1301, USA.
 */

#include "config.h"
#if ENABLE(FILTERS) && ENABLE(OPENCL)
#include "Filter.h"
#include "ImageBuffer.h"
#include "SourceAlpha.h"
#include "OpenCLContext.h"

namespace WebCore {

bool SourceAlpha::platformApplyOpenCL()
{
    FilterContextOpenCL* context = FilterContextOpenCL::context();

    if (!context)
        return false;

    Filter* filter = this->filter();
    if (!filter->sourceImage()) {
        return true;
    }

    ImageBuffer* sourceImage = filter->sourceImage();
    RefPtr<Uint8ClampedArray> sourceImageData = sourceImage->getUnmultipliedImageData(IntRect(IntPoint(),sourceImage->internalSize()));

    unsigned char* dst = sourceImageData->data();
    unsigned char* end = dst + sourceImageData->length();
    while (dst < end) {
        dst[0] = 0;
        dst[1] = 0;
        dst[2] = 0;
        dst += 4;
    }

    createOpenCLImageResult(sourceImageData->data());

    return true;
}

} // namespace WebCore

#endif // ENABLE(FILTERS) && ENABLE(OPENCL)



/*
 * Copyright (C) 2012 Tamas Czene <tczene@inf.u-szeged.hu>
 *
 * This library is free software; you can redistribute it and/or
 * modify it under the terms of the GNU Library General Public
 * License as published by the Free Software Foundation; either
 * version 2 of the License, or (at your option) any later version.
 *
 * This library is distributed in the hope that it will be useful,
 * but WITHOUT ANY WARRANTY; without even the implied warranty of
 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
 * Library General Public License for more details.
 *
 * You should have received a copy of the GNU Library General Public License
 * along with this library; see the file COPYING.LIB.  If not, write to
 * the Free Software Foundation, Inc., 51 Franklin Street, Fifth Floor,
 * Boston, MA 02110-1301, USA.
 */

#include "config.h"
#if ENABLE(FILTERS) && ENABLE(OPENCL)
#include "Filter.h"
#include "ImageBuffer.h"
#include "SourceGraphic.h"
#include "OpenCLContext.h"

namespace WebCore {

bool SourceGraphic::platformApplyOpenCL()
{
    FilterContextOpenCL* context = FilterContextOpenCL::context();

    if (!context)
        return false;

    Filter* filter = this->filter();
    if (!filter->sourceImage()) {
        return true;
    }

    ImageBuffer* sourceImage = filter->sourceImage();
    RefPtr<Uint8ClampedArray> sourceImageData = sourceImage->getPremultipliedImageData(IntRect(IntPoint(),sourceImage->internalSize()));

    createOpenCLImageResult(sourceImageData->data());

    return true;
}

} // namespace WebCore

#endif // ENABLE(FILTERS) && ENABLE(OPENCL)



/*
 * Copyright (C) 2012 Tamas Czene <tczene@inf.u-szeged.hu>
 *
 * This library is free software; you can redistribute it and/or
 * modify it under the terms of the GNU Library General Public
 * License as published by the Free Software Foundation; either
 * version 2 of the License, or (at your option) any later version.
 *
 * This library is distributed in the hope that it will be useful,
 * but WITHOUT ANY WARRANTY; without even the implied warranty of
 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
 * Library General Public License for more details.
 *
 * You should have received a copy of the GNU Library General Public License
 * along with this library; see the file COPYING.LIB.  If not, write to
 * the Free Software Foundation, Inc., 51 Franklin Street, Fifth Floor,
 * Boston, MA 02110-1301, USA.
 */

#include "config.h"
#if ENABLE(FILTERS) && ENABLE(OPENCL)
#include "FETile.h"
#include "OpenCLContext.h"

#include <stdio.h>

#define PROGRAM_STR(Src) #Src
#define PROGRAM(Src) PROGRAM_STR(Src)

namespace WebCore {

static const char* FETileKernelProgram =
PROGRAM_STR(
    __kernel void Tiling(__write_only image2d_t destination,
                     __read_only image2d_t source,
                       float width,
                       float height,
                       float x,
                       float y,
                       sampler_t sampler)
    {
        int coordX = get_global_id(0);
        int coordY = get_global_id(1);

        int2 unnormalizedCoord = (int2)(coordX,coordY);
        float2 normalizedCoord = (float2)((coordX + 0.5f + x) / width, (coordY + 0.5f + y) / height);

        float4 sourcePixel = read_imagef(source, sampler, normalizedCoord);
        write_imagef(destination, unnormalizedCoord, sourcePixel);
    }



__kernel void TileBuilder(__write_only image2d_t destination,
                          __read_only image2d_t source,
                          float sourceWidth,
                          float sourceHeight,
                          sampler_t sampler)
{

    int2 coord = (int2) (get_global_id(0) , get_global_id(1) );
    float4 sourcePixel = (float4)(1,1,1,1);

    if(coord.x <= sourceWidth && coord.y <= sourceHeight)
        sourcePixel = read_imagef(source, sampler, coord);

    write_imagef(destination, coord, sourcePixel);
}

);

inline void FilterContextOpenCL::compileFETile()
{
    if (!m_feTile) {
        m_feTileProgram = compileProgram(FETileKernelProgram);
        ASSERT(m_feTileProgram);

        m_feTile = kernelByName(m_feTileProgram, "Tiling");
        ASSERT(m_feTile);

        m_TileBuilder = kernelByName(m_feTileProgram, "TileBuilder");

    }
}

inline void FilterContextOpenCL::applyFETiling(cl_mem destination, IntSize destinationSize, cl_mem source, FloatRect sourceRect, FloatPoint relativeSourceLocation)
{
    cl_sampler sampler = clCreateSampler(deviceContext(), CL_TRUE, CL_ADDRESS_REPEAT,
                                          CL_FILTER_LINEAR, 0);
    RunKernel kernel(this, m_feTile, destinationSize.width(), destinationSize.height());
    kernel.addArgument(destination);
    kernel.addArgument(source);
    kernel.addArgument(sourceRect.width());
    kernel.addArgument(sourceRect.height());
    kernel.addArgument(relativeSourceLocation.x());
    kernel.addArgument(relativeSourceLocation.y());
    kernel.addArgument(sampler);
    kernel.run();

}

inline void FilterContextOpenCL::applyFETileBuilder(cl_mem destination, FloatRect destinationSize, cl_mem source, FloatSize sourceSize)
{
    cl_sampler sampler = clCreateSampler(deviceContext(), CL_FALSE, CL_ADDRESS_NONE, CL_FILTER_LINEAR, 0);

    RunKernel kernel2(this, m_TileBuilder, destinationSize.width(), destinationSize.height());
    kernel2.addArgument(destination);
    kernel2.addArgument(source);
    kernel2.addArgument(sourceSize.width());
    kernel2.addArgument(sourceSize.height());
    kernel2.addArgument(sampler);
    kernel2.run();

}

bool FETile::platformApplyOpenCL()
{
    FilterContextOpenCL* context = FilterContextOpenCL::context();

    if (!context)
        return false;

    context->compileFETile();

    FilterEffect* in = inputEffect(0);
    cl_mem source = in->asOpenCLImage();
    cl_mem destination = createOpenCLImageResult();

    FloatRect tileRect = in->maxEffectRect();

    if (in->filterEffectType() == FilterEffectTypeSourceInput) {
        Filter* filter = this->filter();
        tileRect = filter->filterRegion();
        tileRect.scale(filter->filterResolution().width(), filter->filterResolution().height());
    }

    cl_mem tile = context->createOpenCLImage(tileRect.size());

    context->applyFETileBuilder(tile, tileRect, source, in->absolutePaintRect().size());

    FloatPoint relativeSourceLocation(in->absolutePaintRect().location());
    relativeSourceLocation.setX(absolutePaintRect().x() - relativeSourceLocation.x());
    relativeSourceLocation.setY(absolutePaintRect().y() - relativeSourceLocation.y());

    context->applyFETiling(destination, absolutePaintRect().size(), tile, tileRect, relativeSourceLocation);

    return true;
}

} // namespace WebCore

#endif // ENABLE(FILTERS) && ENABLE(OPENCL)



/*
 * Copyright (C) 2012 Tamas Czene <tczene@inf.u-szeged.hu>
 *
 * This library is free software; you can redistribute it and/or
 * modify it under the terms of the GNU Library General Public
 * License as published by the Free Software Foundation; either
 * version 2 of the License, or (at your option) any later version.
 *
 * This library is distributed in the hope that it will be useful,
 * but WITHOUT ANY WARRANTY; without even the implied warranty of
 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
 * Library General Public License for more details.
 *
 * You should have received a copy of the GNU Library General Public License
 * along with this library; see the file COPYING.LIB.  If not, write to
 * the Free Software Foundation, Inc., 51 Franklin Street, Fifth Floor,
 * Boston, MA 02110-1301, USA.
 */

#include "config.h"
#if ENABLE(FILTERS) && ENABLE(OPENCL)
#include "FETurbulence.h"
#include "FETurbulence.cpp"
#include "OpenCLContext.h"
#include "SVGFilter.h"
#include <fstream>
#include <stdio.h>

#define PROGRAM_STR(Src) #Src
#define PROGRAM(Src) PROGRAM_STR(Src)

namespace WebCore {

static const char* FETurbulenceKernelProgram =
PROGRAM_STR(
__constant int s_perlinNoise = 4096;
__constant int s_blockSize = 256;
__constant int s_blockMask = 255;

typedef struct
{
    int noisePositionIntegerValue;
    float noisePositionFractionValue;
} Noise;

typedef struct
{
    int width;
    int wrapX;
    int height;
    int wrapY;
} StitchData;

float linearInterpolation(float t, float a, float b)
{
    return a + t * (b - a);
}

float noise2D(__constant float *component, __constant int *latticeSelector, StitchData stitchData, float noiseVectorX, float noiseVectorY, int stitchTiles)
{
    Noise noiseX;
    noiseX.noisePositionIntegerValue = (int)(noiseVectorX + s_perlinNoise);
    noiseX.noisePositionFractionValue = (noiseVectorX + s_perlinNoise) - noiseX.noisePositionIntegerValue;
    Noise noiseY;
    noiseY.noisePositionIntegerValue = (int)(noiseVectorY + s_perlinNoise);
    noiseY.noisePositionFractionValue = (noiseVectorY + s_perlinNoise) - noiseY.noisePositionIntegerValue;

    // If stitching, adjust lattice points accordingly.
    if (stitchTiles) {
        if (noiseX.noisePositionIntegerValue >= stitchData.wrapX)
            noiseX.noisePositionIntegerValue -= stitchData.width;
        if (noiseX.noisePositionIntegerValue >= stitchData.wrapX - 1)
            noiseX.noisePositionIntegerValue -= stitchData.width - 1;
        if (noiseY.noisePositionIntegerValue >= stitchData.wrapY)
            noiseY.noisePositionIntegerValue -= stitchData.height;
        if (noiseY.noisePositionIntegerValue >= stitchData.wrapY - 1)
            noiseY.noisePositionIntegerValue -= stitchData.height - 1;
    }

    noiseX.noisePositionIntegerValue &= s_blockMask;
    noiseY.noisePositionIntegerValue &= s_blockMask;
    int latticeIndex = latticeSelector[noiseX.noisePositionIntegerValue];
    int nextLatticeIndex = latticeSelector[(noiseX.noisePositionIntegerValue + 1) & s_blockMask];

    float sx = noiseX.noisePositionFractionValue * noiseX.noisePositionFractionValue * (3 - 2 * noiseX.noisePositionFractionValue);
    float sy = noiseY.noisePositionFractionValue * noiseY.noisePositionFractionValue * (3 - 2 * noiseY.noisePositionFractionValue);

    // This is taken 1:1 from SVG spec: http://www.w3.org/TR/SVG11/filters.html#feTurbulenceElement.
    int temp = latticeSelector[latticeIndex + noiseY.noisePositionIntegerValue];
    float u = noiseX.noisePositionFractionValue * component[temp * 2] + noiseY.noisePositionFractionValue * component[temp * 2 + 1];
    temp = latticeSelector[nextLatticeIndex + noiseY.noisePositionIntegerValue];
    float v = (noiseX.noisePositionFractionValue - 1) * component[temp * 2] + noiseY.noisePositionFractionValue * component[temp * 2 + 1];
    float a = linearInterpolation(sx, u, v);
    temp = latticeSelector[latticeIndex + noiseY.noisePositionIntegerValue + 1];
    u = noiseX.noisePositionFractionValue * component[temp * 2] + (noiseY.noisePositionFractionValue - 1) * component[temp * 2 + 1];
    temp = latticeSelector[nextLatticeIndex + noiseY.noisePositionIntegerValue + 1];
    v = (noiseX.noisePositionFractionValue - 1) * component[temp * 2] + (noiseY.noisePositionFractionValue - 1) * component[temp * 2 + 1];
    float b = linearInterpolation(sx, u, v);
    return linearInterpolation(sy, a, b);
}

__kernel void Turbulence(__write_only image2d_t destination,
                         __constant float *transform,
                         __constant float *redComponent,
                         __constant float *greenComponent,
                         __constant float *blueComponent,
                         __constant float *alphaComponent,
                         __constant int *latticeSelector,
                         __private int offsetX,
                         __private int offsetY,
                         __private int tileWidth,
                         __private int tileHeight,
                         __private float baseFrequencyX,
                         __private float baseFrequencyY,
                         __private int stitchTiles,
                         __private int numOctaves,
                         __private int type,
                         __private int filter_height)
{
    StitchData stitchData;
    stitchData.width = 0;
    stitchData.wrapX = 0;
    stitchData.height = 0;
    stitchData.wrapY = 0;
    // Adjust the base frequencies if necessary for stitching.
    if (stitchTiles) {
        // When stitching tiled turbulence, the frequencies must be adjusted
        // so that the tile borders will be continuous.
        if (baseFrequencyX) {
            float lowFrequency = floor(tileWidth * baseFrequencyX) / tileWidth;
            float highFrequency = ceil(tileWidth * baseFrequencyX) / tileWidth;
            // BaseFrequency should be non-negative according to the standard.
            if (baseFrequencyX / lowFrequency < highFrequency / baseFrequencyX)
                baseFrequencyX = lowFrequency;
            else
                baseFrequencyX = highFrequency;
        }
        if (baseFrequencyY) {
            float lowFrequency = floor(tileHeight * baseFrequencyY) / tileHeight;
            float highFrequency = ceil(tileHeight * baseFrequencyY) / tileHeight;
            if (baseFrequencyY / lowFrequency < highFrequency / baseFrequencyY)
                baseFrequencyY = lowFrequency;
            else
                baseFrequencyY = highFrequency;
        }
        // Set up TurbulenceInitial stitch values.
        stitchData.width = round(tileWidth * baseFrequencyX);
        stitchData.wrapX = s_perlinNoise + stitchData.width;
        stitchData.height = round(tileHeight * baseFrequencyY);
        stitchData.wrapY = s_perlinNoise + stitchData.height;
    }
    float4 turbulenceFunctionResult = (float4)(0,0,0,0);
    float x = (get_global_id(0) + offsetX) * baseFrequencyX;
    float y = (get_global_id(1) + offsetY) * baseFrequencyY;

    float noiseVectorX = transform[0] * x + transform[2] * y + transform[4];
    float noiseVectorY = transform[1] * x + transform[3] * y + transform[5];

    float ratio = 1;
    for (int octave = 0; octave < numOctaves; ++octave) {
        if (type == 1) {
            turbulenceFunctionResult.x += noise2D(redComponent, latticeSelector, stitchData, noiseVectorX, noiseVectorY, stitchTiles) / ratio;
            turbulenceFunctionResult.y += noise2D(greenComponent, latticeSelector, stitchData, noiseVectorX, noiseVectorY, stitchTiles) / ratio;
            turbulenceFunctionResult.z += noise2D(blueComponent, latticeSelector, stitchData, noiseVectorX, noiseVectorY, stitchTiles) / ratio;
            turbulenceFunctionResult.w += noise2D(alphaComponent, latticeSelector, stitchData, noiseVectorX, noiseVectorY, stitchTiles) / ratio;
        } else {
            turbulenceFunctionResult.x += fabs(noise2D(redComponent, latticeSelector, stitchData, noiseVectorX, noiseVectorY, stitchTiles)) / ratio;
            turbulenceFunctionResult.y += fabs(noise2D(greenComponent, latticeSelector, stitchData, noiseVectorX, noiseVectorY, stitchTiles)) / ratio;
            turbulenceFunctionResult.z += fabs(noise2D(blueComponent, latticeSelector, stitchData, noiseVectorX, noiseVectorY, stitchTiles)) / ratio;
            turbulenceFunctionResult.w += fabs(noise2D(alphaComponent, latticeSelector, stitchData, noiseVectorX, noiseVectorY, stitchTiles)) / ratio;
        }
        noiseVectorX *= 2;
        noiseVectorY *= 2;
        ratio *= 2;
        if (stitchTiles) {
            // Update stitch values. Subtracting s_perlinNoiseoise before the multiplication and
            // adding it afterward simplifies to subtracting it once.
            stitchData.width *= 2;
            stitchData.wrapX = 2 * stitchData.wrapX - s_perlinNoise;
            stitchData.height *= 2;
            stitchData.wrapY = 2 * stitchData.wrapY - s_perlinNoise;
        }
    }

    if (type == 1)
        turbulenceFunctionResult = turbulenceFunctionResult * 0.5f + 0.5f;
    // Clamp result
    turbulenceFunctionResult = clamp(turbulenceFunctionResult, 0.0f, 1.0f);

    write_imagef(destination, (int2)(get_global_id(0), get_global_id(1)), turbulenceFunctionResult);
}
);

inline void FilterContextOpenCL::compileFETurbulence()
{

    if (!m_FETurbulence) {
        m_FETurbulenceProgram = compileProgram(FETurbulenceKernelProgram);
        ASSERT(m_FETurbulenceProgram);
        m_FETurbulence = kernelByName(m_FETurbulenceProgram, "Turbulence");
        ASSERT(m_FETurbulence);
    }
}

inline void FilterContextOpenCL::applyFETurbulence(cl_mem destination, IntSize destinationSize, cl_mem transformHandle, cl_mem redComponentHandle, cl_mem greenComponentHandle,
                                                   cl_mem blueComponentHandle, cl_mem alphaComponentHandle, cl_mem latticeSelectorHandle, int offsetX, int offsetY, int tileWidth, int tileHeight,
                                                   float baseFrequencyX, float baseFrequencyY, bool stitchTiles, int numOctaves, int type)
{
    RunKernel kernel(this, m_FETurbulence, destinationSize.width(), destinationSize.height());
    kernel.addArgument(destination);
    kernel.addArgument(transformHandle);
    kernel.addArgument(redComponentHandle);
    kernel.addArgument(greenComponentHandle);
    kernel.addArgument(blueComponentHandle);
    kernel.addArgument(alphaComponentHandle);
    kernel.addArgument(latticeSelectorHandle);
    kernel.addArgument(offsetX);
    kernel.addArgument(offsetY);
    kernel.addArgument(tileWidth);
    kernel.addArgument(tileHeight);
    kernel.addArgument(baseFrequencyX);
    kernel.addArgument(baseFrequencyY);
    kernel.addArgument(stitchTiles);
    kernel.addArgument(numOctaves);
    kernel.addArgument(type);
    kernel.addArgument(destinationSize.height());
    kernel.run();
}

bool FETurbulence::platformApplyOpenCL()
{
    FilterContextOpenCL* context = FilterContextOpenCL::context();
    if (!context)
        return false;

    context->compileFETurbulence();

    cl_mem destination = createOpenCLImageResult();

    PaintingData paintingData(m_seed, roundedIntSize(filterPrimitiveSubregion().size()));
    initPaint(paintingData);

    AffineTransform invertedTransform = reinterpret_cast<SVGFilter*>(filter())->absoluteTransform().inverse();
    float transformComponents[6] = { invertedTransform.a(), invertedTransform.b(), invertedTransform.c(), invertedTransform.d(), invertedTransform.e(), invertedTransform.f() };

    cl_mem transformHandle = context->uploadBuffer(context, transformComponents, sizeof(float) * 6);
    cl_mem redComponentHandle = context->uploadBuffer(context, paintingData.gradient + 0, sizeof(float) * (2 * s_blockSize + 2) * 2);
    cl_mem greenComponentHandle = context->uploadBuffer(context, paintingData.gradient + 1, sizeof(float) * (2 * s_blockSize + 2) * 2);
    cl_mem blueComponentHandle = context->uploadBuffer(context, paintingData.gradient + 2, sizeof(float) * (2 * s_blockSize + 2) * 2);
    cl_mem alphaComponentHandle = context->uploadBuffer(context, paintingData.gradient + 3, sizeof(float) * (2 * s_blockSize + 2) * 2);
    cl_mem latticeSelectorHandle = context->uploadBuffer(context, paintingData.latticeSelector, sizeof(int) * (2 * s_blockSize + 2));

    context->applyFETurbulence(destination, absolutePaintRect().size(), transformHandle, redComponentHandle, greenComponentHandle,
                               blueComponentHandle, alphaComponentHandle, latticeSelectorHandle, absolutePaintRect().x(), absolutePaintRect().y(),
                               paintingData.filterSize.width(), paintingData.filterSize.height(), m_baseFrequencyX, m_baseFrequencyY, m_stitchTiles,
                               m_numOctaves, m_type);

    clReleaseMemObject(transformHandle);
    clReleaseMemObject(redComponentHandle);
    clReleaseMemObject(greenComponentHandle);
    clReleaseMemObject(blueComponentHandle);
    clReleaseMemObject(alphaComponentHandle);
    clReleaseMemObject(latticeSelectorHandle);

    return true;
}

} // namespace WebCore

#endif // ENABLE(FILTERS) && ENABLE(OPENCL)



/*
 * Copyright (C) 2012 Tamas Czene <tczene@inf.u-szeged.hu>
 *
 * This library is free software; you can redistribute it and/or
 * modify it under the terms of the GNU Library General Public
 * License as published by the Free Software Foundation; either
 * version 2 of the License, or (at your option) any later version.
 *
 * This library is distributed in the hope that it will be useful,
 * but WITHOUT ANY WARRANTY; without even the implied warranty of
 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
 * Library General Public License for more details.
 *
 * You should have received a copy of the GNU Library General Public License
 * along with this library; see the file COPYING.LIB.  If not, write to
 * the Free Software Foundation, Inc., 51 Franklin Street, Fifth Floor,
 * Boston, MA 02110-1301, USA.
 */

#include "config.h"
#if ENABLE(FILTERS) && ENABLE(OPENCL)
#include "FEOffset.h"
#include "OpenCLContext.h"
#include "stdio.h"

#define PROGRAM_STR(Src) #Src
#define PROGRAM(Src) PROGRAM_STR(Src)

struct BMPHeader {
  unsigned short bfType;
  unsigned int bfSize;
  unsigned short bfReserved1;
  unsigned short bfReserved2;
  unsigned int bfOffBits;

  unsigned int biSize;
  unsigned int biWidth;
  unsigned int biHeight;
  unsigned short biPlanes;
  unsigned short biBitCount;
  unsigned int biCompression;
  unsigned int biSizeImage;
  unsigned int biXPelsPerMeter;
  unsigned int biYPelsPerMeter;
  unsigned int biClrUsed;
  unsigned int biClrImportant;
 } __attribute__((__packed__));

 static void dumpAsBMP(const char* fileName, RefPtr<Uint8ClampedArray> pixels, int x, int y, int width, int height, int srcScanline)
 {
  BMPHeader* header;
  unsigned char* src;
  unsigned char* dst;
  int dstScanline;
  int size, i, j;
  FILE* f;

  dstScanline = ((width * 3) + 3) & ~3;
  srcScanline *= 4;

  size = (dstScanline * height) + sizeof(BMPHeader);
  header = (BMPHeader*)fastMalloc(size);
  if (!header)
  return;

  f = fopen(fileName, "w");
  if (!f) {
  fastFree(header);
  return;
  }

  memset(header, 0, size);
  header->bfType = 0x4d42;
  header->bfSize = size;
  header->bfOffBits = sizeof(BMPHeader);
  header->biSize = 40;
  header->biWidth = width;
  header->biHeight = height;
  header->biPlanes = 1;
  header->biBitCount = 24;
  size -= sizeof(BMPHeader);
  header->biSizeImage = size;

  dst = ((unsigned char*)header) + sizeof(BMPHeader);
  src = pixels->data() + ((y + height - 1) * srcScanline) + x * 4;

  dstScanline -= (width * 3);
  srcScanline += (width * 4);
  for (i = 0; i < height; ++i) {
  for (j = 0; j < width; ++j) {
  dst[0] = src[2];
  dst[1] = src[1];
  dst[2] = src[0];
  src += 4;
  dst += 3;
  }
  dst += dstScanline;
  src -= srcScanline;
  }

  fwrite(header, size, 1, f);
  fclose(f);
  fastFree(header);
  printf("'%s' written\n", fileName);
  fflush(stdout);
}

namespace WebCore {

static const char* FEOffsetKernelProgram =
PROGRAM_STR(
    __kernel void Offset(__write_only image2d_t destination,
                     __read_only image2d_t source,
                      float x,
                      float y,
                      sampler_t sampler)
    {
        int2 sourceCoord = (int2) (get_global_id(0) + x  , get_global_id(1) + y);
        float4 sourcePixel = read_imagef(source, sampler, sourceCoord);
        write_imagef(destination, (int2)(get_global_id(0), get_global_id(1)), sourcePixel);
    }

);

inline void FilterContextOpenCL::compileFEOffset()
{
    if (!m_feOffset) {
        m_feOffsetProgram = compileProgram(FEOffsetKernelProgram);
        ASSERT(m_feOffsetProgram);
        m_feOffset = kernelByName(m_feOffsetProgram, "Offset");
        ASSERT(m_feOffset);
    }
}

inline void FilterContextOpenCL::applyFEOffset(cl_mem destination, IntSize destinationSize, cl_mem source, FloatPoint relativeSourceLocation)
{

    float x = relativeSourceLocation.x();
    float y = relativeSourceLocation.y();

    cl_sampler sampler = clCreateSampler(deviceContext(), CL_FALSE, CL_ADDRESS_CLAMP,
                                         CL_FILTER_LINEAR, 0);

    RunKernel kernel(this, m_feOffset, destinationSize.width(), destinationSize.height());
    kernel.addArgument(destination);
    kernel.addArgument(source);
    kernel.addArgument(x);
    kernel.addArgument(y);
    kernel.addArgument(sampler);
    kernel.run();
}

bool FEOffset::platformApplyOpenCL()
{
    FilterContextOpenCL* context = FilterContextOpenCL::context();

    if (!context)
        return false;

    context->compileFEOffset();

    FilterEffect* in = inputEffect(0);
    cl_mem source = in->asOpenCLImage();
    cl_mem destination = createOpenCLImageResult();

    size_t origin[3] = { 0, 0, 0 };
    size_t region[3] = { absolutePaintRect().width(), absolutePaintRect().height(), 1};

    RefPtr<Uint8ClampedArray> destinationPixelArray = Uint8ClampedArray::create(absolutePaintRect().width() * absolutePaintRect().height() * 4);

    clEnqueueReadImage(context->commandQueue(), source, CL_TRUE, origin, region, 0, 0, destinationPixelArray->data(), 0, 0, 0);

    dumpAsBMP("kep3.bmp", destinationPixelArray, 0, 0, absolutePaintRect().width(), absolutePaintRect().height(), absolutePaintRect().width());

    clEnqueueReadImage(context->commandQueue(), destination, CL_TRUE, origin, region, 0, 0, destinationPixelArray->data(), 0, 0, 0);

    dumpAsBMP("kep4.bmp", destinationPixelArray, 0, 0, absolutePaintRect().width(), absolutePaintRect().height(), absolutePaintRect().width());


    FloatPoint relativeSourceLocation(in->absolutePaintRect().location());
    //relativeSourceLocation.setX(absolutePaintRect().x() - (relativeSourceLocation.x() + filter()->applyHorizontalScale(m_dx)));
    //relativeSourceLocation.setY(absolutePaintRect().y() - (relativeSourceLocation.y() + filter()->applyVerticalScale(m_dy)));
    context->applyFEOffset(destination, absolutePaintRect().size(), source,relativeSourceLocation);

    clEnqueueReadImage(context->commandQueue(), destination, CL_TRUE, origin, region, 0, 0, destinationPixelArray->data(), 0, 0, 0);

    dumpAsBMP("kep5.bmp", destinationPixelArray, 0, 0, absolutePaintRect().width(), absolutePaintRect().height(), absolutePaintRect().width());

   // platformApplySoftware();



    dumpAsBMP("kep7.bmp", asImageBuffer()->getPremultipliedImageData(absolutePaintRect()), 0, 0, absolutePaintRect().width(), absolutePaintRect().height(), absolutePaintRect().width());


    return true;
}

} // namespace WebCore

#endif // ENABLE(FILTERS) && ENABLE(OPENCL)

Lines 33-38 public: a/Source/WebCore/platform/graphics/filters/SourceAlpha.h_sec1
- a/Source/WebCore/platform/graphics/filters/SourceAlpha.h +3 lines
33	static const AtomicString& effectName();	33	static const AtomicString& effectName();
34		34
35	virtual void platformApplySoftware();	35	virtual void platformApplySoftware();
		36	#if ENABLE(OPENCL)
		37	virtual bool platformApplyOpenCL();
		38	#endif
36	virtual void dump();	39	virtual void dump();
37		40
38	virtual void determineAbsolutePaintRect();	41	virtual void determineAbsolutePaintRect();

Lines 34-39 public: a/Source/WebCore/platform/graphics/filters/SourceGraphic.h_sec1
- a/Source/WebCore/platform/graphics/filters/SourceGraphic.h +3 lines
34	static const AtomicString& effectName();	34	static const AtomicString& effectName();
35		35
36	virtual void platformApplySoftware();	36	virtual void platformApplySoftware();
		37	#if ENABLE(OPENCL)
		38	virtual bool platformApplyOpenCL();
		39	#endif
37	virtual void dump();	40	virtual void dump();
38		41
39	virtual void determineAbsolutePaintRect();	42	virtual void determineAbsolutePaintRect();



            lastEffect->apply();
            lastEffect->correctFilterResultIfNeeded();
#if !USE(CG)
#if ENABLE(OPENCL)
        cl_mem resultImage = lastEffect->asOpenCLImage();
        FilterContextOpenCL* context = FilterContextOpenCL::context();
        if (resultImage && context)
            context->OpenCLTransformColorSpace(resultImage, ColorSpaceDeviceRGB, lastEffect->colorSpace(), lastEffect->absolutePaintRect());
#else
        ImageBuffer* resultImage = lastEffect->asImageBuffer();
        if (resultImage)
            resultImage->transformColorSpace(lastEffect->colorSpace(), ColorSpaceDeviceRGB);
#endif
#endif
        }
        filterData->builded = true;

Lines 39-44 public: a/Source/WebCore/svg/graphics/filters/SVGFEImage.h_sec1
- a/Source/WebCore/svg/graphics/filters/SVGFEImage.h +3 lines
39	static PassRefPtr<FEImage> createWithIRIReference(Filter, Document, const String&, const SVGPreserveAspectRatio&);	39	static PassRefPtr<FEImage> createWithIRIReference(Filter, Document, const String&, const SVGPreserveAspectRatio&);
40		40
41	virtual void platformApplySoftware();	41	virtual void platformApplySoftware();
		42	#if ENABLE(OPENCL)
		43	virtual bool platformApplyOpenCL();
		44	#endif
42	virtual void dump();	45	virtual void dump();
43		46
44	virtual void determineAbsolutePaintRect();	47	virtual void determineAbsolutePaintRect();

Return to Bug 70099

Lines 4014-4019 contains(CONFIG, opengl-shims) { a/Source/WebCore/Target.pri_sec1
- a/Source/WebCore/Target.pri +20 lines
4014	DEFINES += QT_OPENGL_SHIMS=1	4014	DEFINES += QT_OPENGL_SHIMS=1
4015	}	4015	}
4016		4016
		4017	contains(DEFINES, ENABLE_OPENCL=1) {
		4018	HEADERS += platform/graphics/filters/OpenCL/OpenCLContext.h
		4019	SOURCES += platform/graphics/filters/OpenCL/OpenCLContext.cpp
		4020	SOURCES += platform/graphics/filters/OpenCL/OpenCLFEImage.cpp
		4021	SOURCES += platform/graphics/filters/OpenCL/OpenCLFEComposite.cpp
		4022	SOURCES += platform/graphics/filters/OpenCL/OpenCLFEBlend.cpp
		4023	SOURCES += platform/graphics/filters/OpenCL/OpenCLFEFlood.cpp
		4024	SOURCES += platform/graphics/filters/OpenCL/OpenCLFEOffset.cpp
		4025	SOURCES += platform/graphics/filters/OpenCL/OpenCLFESourceGraphic.cpp
		4026	SOURCES += platform/graphics/filters/OpenCL/OpenCLFESourceAlpha.cpp
		4027	SOURCES += platform/graphics/filters/OpenCL/OpenCLFETurbulence.cpp
		4028	SOURCES += platform/graphics/filters/OpenCL/OpenCLFEColorMatrix.cpp
		4029	SOURCES += platform/graphics/filters/OpenCL/OpenCLFEConvolveMatrix.cpp
		4030	#SOURCES += platform/graphics/filters/OpenCL/OpenCLFEComponentTransfer.cpp
		4031	SOURCES += platform/graphics/filters/OpenCL/OpenCLFEMerge.cpp
		4032	SOURCES += platform/graphics/filters/OpenCL/OpenCLFEMorphology.cpp
		4033	SOURCES += platform/graphics/filters/OpenCL/OpenCLFETile.cpp
		4034	DEFINES += OPENCL=1
		4035	}
		4036
4017	use?(GRAPHICS_SURFACE) {	4037	use?(GRAPHICS_SURFACE) {
4018	mac {	4038	mac {
4019	SOURCES += platform/graphics/surfaces/mac/GraphicsSurfaceMac.cpp	4039	SOURCES += platform/graphics/surfaces/mac/GraphicsSurfaceMac.cpp

Lines 305-310 void FEGaussianBlur::platformApplySoftware() a/Source/WebCore/platform/graphics/filters/FEGaussianBlur.cpp_sec1
- a/Source/WebCore/platform/graphics/filters/FEGaussianBlur.cpp +12 lines
305	platformApply(srcPixelArray, tmpPixelArray, kernelSizeX, kernelSizeY, paintSize);	305	platformApply(srcPixelArray, tmpPixelArray, kernelSizeX, kernelSizeY, paintSize);
306	}	306	}
307		307
		308	#if ENABLE(OPENCL)
		309	bool FEGaussianBlur::platformApplyOpenCL()
		310	{
		311	asImageBuffer();
		312	platformApplySoftware();
		313	ImageBuffer* sourceImage = asImageBuffer();
		314	RefPtr<Uint8ClampedArray> sourceImageData = sourceImage->getUnmultipliedImageData(IntRect(IntPoint(),sourceImage->internalSize()));
		315	createOpenCLImageResult(sourceImageData->data());
		316	return true;
		317	}
		318	#endif
		319
308	void FEGaussianBlur::dump()	320	void FEGaussianBlur::dump()
309	{	321	{
310	}	322	}

Lines 302-310 void RenderSVGResourceFilter::postApplyResource(RenderObject* object, GraphicsCo a/Source/WebCore/rendering/svg/RenderSVGResourceFilter.cpp_sec1
- a/Source/WebCore/rendering/svg/RenderSVGResourceFilter.cpp -3 / +10 lines
302	lastEffect->apply();	302	lastEffect->apply();
303	lastEffect->correctFilterResultIfNeeded();	303	lastEffect->correctFilterResultIfNeeded();
304	#if !USE(CG)	304	#if !USE(CG)
305	ImageBuffer* resultImage = lastEffect->asImageBuffer();	305	#if ENABLE(OPENCL)
306	if (resultImage)	306	cl_mem resultImage = lastEffect->asOpenCLImage();
307	resultImage->transformColorSpace(lastEffect->colorSpace(), ColorSpaceDeviceRGB);	307	FilterContextOpenCL* context = FilterContextOpenCL::context();
		308	if (resultImage && context)
		309	context->OpenCLTransformColorSpace(resultImage, ColorSpaceDeviceRGB, lastEffect->colorSpace(), lastEffect->absolutePaintRect());
		310	#else
		311	ImageBuffer* resultImage = lastEffect->asImageBuffer();
		312	if (resultImage)
		313	resultImage->transformColorSpace(lastEffect->colorSpace(), ColorSpaceDeviceRGB);
		314	#endif
308	#endif	315	#endif
309	}	316	}
310	filterData->builded = true;	317	filterData->builded = true;

Lines 174-179 void FEComponentTransfer::platformApplySoftware() a/Source/WebCore/platform/graphics/filters/FEComponentTransfer.cpp_sec1
- a/Source/WebCore/platform/graphics/filters/FEComponentTransfer.cpp +12 lines
174	}	174	}
175	}	175	}
176		176
		177	#if ENABLE(OPENCL)
		178	bool FEComponentTransfer::platformApplyOpenCL()
		179	{
		180	asImageBuffer();
		181	platformApplySoftware();
		182	ImageBuffer* sourceImage = asImageBuffer();
		183	RefPtr<Uint8ClampedArray> sourceImageData = sourceImage->getUnmultipliedImageData(IntRect(IntPoint(),sourceImage->internalSize()));
		184	createOpenCLImageResult(sourceImageData->data());
		185	return true;
		186	}
		187	#endif
		188
177	void FEComponentTransfer::getValues(unsigned char rValues[256], unsigned char gValues[256], unsigned char bValues[256], unsigned char aValues[256])	189	void FEComponentTransfer::getValues(unsigned char rValues[256], unsigned char gValues[256], unsigned char bValues[256], unsigned char aValues[256])
178	{	190	{
179	for (unsigned i = 0; i < 256; ++i)	191	for (unsigned i = 0; i < 256; ++i)

Lines 135-140 void FECustomFilter::platformApplySoftware() a/Source/WebCore/platform/graphics/filters/FECustomFilter.cpp_sec1
- a/Source/WebCore/platform/graphics/filters/FECustomFilter.cpp +12 lines
135	clearShaderResult();	135	clearShaderResult();
136	}	136	}
137		137
		138	#if ENABLE(OPENCL)
		139	bool FECustomFilter::platformApplyOpenCL()
		140	{
		141	asImageBuffer();
		142	platformApplySoftware();
		143	ImageBuffer* sourceImage = asImageBuffer();
		144	RefPtr<Uint8ClampedArray> sourceImageData = sourceImage->getUnmultipliedImageData(IntRect(IntPoint(),sourceImage->internalSize()));
		145	createOpenCLImageResult(sourceImageData->data());
		146	return true;
		147	}
		148	#endif
		149
138	void FECustomFilter::clearShaderResult()	150	void FECustomFilter::clearShaderResult()
139	{	151	{
140	clearResult();	152	clearResult();

Lines 133-138 void FEDisplacementMap::platformApplySoftware() a/Source/WebCore/platform/graphics/filters/FEDisplacementMap.cpp_sec1
- a/Source/WebCore/platform/graphics/filters/FEDisplacementMap.cpp +12 lines
133	}	133	}
134	}	134	}
135		135
		136	#if ENABLE(OPENCL)
		137	bool FEDisplacementMap::platformApplyOpenCL()
		138	{
		139	asImageBuffer();
		140	platformApplySoftware();
		141	ImageBuffer* sourceImage = asImageBuffer();
		142	RefPtr<Uint8ClampedArray> sourceImageData = sourceImage->getUnmultipliedImageData(IntRect(IntPoint(),sourceImage->internalSize()));
		143	createOpenCLImageResult(sourceImageData->data());
		144	return true;
		145	}
		146	#endif
		147
136	void FEDisplacementMap::dump()	148	void FEDisplacementMap::dump()
137	{	149	{
138	}	150	}

Lines 119-124 void FEDropShadow::platformApplySoftware() a/Source/WebCore/platform/graphics/filters/FEDropShadow.cpp_sec1
- a/Source/WebCore/platform/graphics/filters/FEDropShadow.cpp +12 lines
119	resultImage->context()->drawImageBuffer(sourceImage, ColorSpaceDeviceRGB, drawingRegion);	119	resultImage->context()->drawImageBuffer(sourceImage, ColorSpaceDeviceRGB, drawingRegion);
120	}	120	}
121		121
		122	#if ENABLE(OPENCL)
		123	bool FEDropShadow::platformApplyOpenCL()
		124	{
		125	asImageBuffer();
		126	platformApplySoftware();
		127	ImageBuffer* sourceImage = asImageBuffer();
		128	RefPtr<Uint8ClampedArray> sourceImageData = sourceImage->getUnmultipliedImageData(IntRect(IntPoint(),sourceImage->internalSize()));
		129	createOpenCLImageResult(sourceImageData->data());
		130	return true;
		131	}
		132	#endif
		133
122	void FEDropShadow::dump()	134	void FEDropShadow::dump()
123	{	135	{
124	}	136	}

Lines 406-411 void FELighting::platformApplySoftware() a/Source/WebCore/platform/graphics/filters/FELighting.cpp_sec1
- a/Source/WebCore/platform/graphics/filters/FELighting.cpp +12 lines
406	drawLighting(srcPixelArray, absolutePaintSize.width(), absolutePaintSize.height());	406	drawLighting(srcPixelArray, absolutePaintSize.width(), absolutePaintSize.height());
407	}	407	}
408		408
		409	#if ENABLE(OPENCL)
		410	bool FELighting::platformApplyOpenCL()
		411	{
		412	asImageBuffer();
		413	platformApplySoftware();
		414	ImageBuffer* sourceImage = asImageBuffer();
		415	RefPtr<Uint8ClampedArray> sourceImageData = sourceImage->getUnmultipliedImageData(IntRect(IntPoint(),sourceImage->internalSize()));
		416	createOpenCLImageResult(sourceImageData->data());
		417	return true;
		418	}
		419	#endif
		420
409	} // namespace WebCore	421	} // namespace WebCore
410		422
411	#endif // ENABLE(FILTERS)	423	#endif // ENABLE(FILTERS)

Line 0 a/Source/WebCore/platform/graphics/filters/OpenCL/OpenCLContext.cpp_sec1
- a/Source/WebCore/platform/graphics/filters/OpenCL/OpenCLContext.cpp +418 lines
		1	#include "config.h"
		2
		3	#if ENABLE(OPENCL)
		4	#include "OpenCLContext.h"
		5	#include <iostream>
		6	#include <time.h>
		7	#include "stdio.h"
		8
		9	#define PROGRAM_STR(Src) #Src
		10	#define PROGRAM(Src) PROGRAM_STR(Src)
		11
		12	namespace WebCore {
		13
		14	FilterContextOpenCL* FilterContextOpenCL::context()
		15	{
		16	static bool wasInitialized = false;
		17	static FilterContextOpenCL* context = 0;
		18
		19	if (wasInitialized)
		20	return context;
		21
		22	#ifdef NDEBUG
		23	printf("debug\n");
		24	#endif
		25	wasInitialized = true;
		26	context = new FilterContextOpenCL();
		27
		28	// Initializing the context.
		29	cl_int errNum;
		30	cl_device_id *devices;
		31	cl_platform_id firstPlatformId;
		32	size_t deviceBufferSize = -1;
		33
		34	errNum = clGetPlatformIDs(1, &firstPlatformId, NULL);
		35	cl_context_properties contextProperties[] = { CL_CONTEXT_PLATFORM, (cl_context_properties)firstPlatformId, 0};
		36	context->m_deviceContext = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_GPU, NULL, NULL, &errNum);
		37	if (errNum != CL_SUCCESS) {
		38	context->m_deviceContext = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_CPU, NULL, NULL, &errNum);
		39	if (errNum != CL_SUCCESS)
		40	return context;
		41	}
		42
		43	errNum = clGetContextInfo(context->m_deviceContext, CL_CONTEXT_DEVICES, 0, NULL, &deviceBufferSize);
		44	if (errNum != CL_SUCCESS)
		45	return context;
		46
		47	if (deviceBufferSize <= 0)
		48	return context;
		49
		50	devices = new cl_device_id[deviceBufferSize / sizeof(cl_device_id)];
		51	errNum = clGetContextInfo(context->m_deviceContext, CL_CONTEXT_DEVICES, deviceBufferSize, devices, NULL);
		52	if (errNum != CL_SUCCESS)
		53	return context;
		54
		55	context->m_commandQueue = clCreateCommandQueue(context->m_deviceContext, devices[0], 0, NULL);
		56	if (context->m_commandQueue == NULL)
		57	return context;
		58
		59	context->m_device = devices[0];
		60	delete [] devices;
		61
		62	cl_bool imageSupport = CL_FALSE;
		63	clGetDeviceInfo(context->m_device, CL_DEVICE_IMAGE_SUPPORT, sizeof(cl_bool), &imageSupport, NULL);
		64	if (imageSupport != CL_TRUE)
65	return context;
66
67	return context;
68	}
69
70	cl_program FilterContextOpenCL::compileProgram(const char* source)
71	{
72	cl_int errNum;
73	cl_program program;
74
75	FilterContextOpenCL* context = FilterContextOpenCL::context();
76
77	program = clCreateProgramWithSource(context->m_deviceContext, 1, (const char**) &source, NULL, &errNum);
78	if (errNum != CL_SUCCESS)
79	OpenCLPrintError(errNum);
80	if (errNum != CL_SUCCESS)
81	return 0;
82
83	errNum = clBuildProgram(program, 0, 0, 0, 0, 0);
84	if (errNum != CL_SUCCESS)
85	OpenCLPrintError(errNum);
86	if (errNum != CL_SUCCESS)
87	return 0;
88
89	return program;
90	}
91
92	static const char* OpenCLTransformColorSpaceKernelProgram =
93	PROGRAM_STR(
94	const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE \| CLK_ADDRESS_CLAMP_TO_EDGE \| CLK_FILTER_NEAREST;
95
96	__kernel void OpenCLTransformColorSpace(__read_only image2d_t source,
97	__write_only image2d_t destination,
98	__constant float *clLookUpTable)
99	{
100	int2 sourceCoord = (int2) (get_global_id(0), get_global_id(1));
101	float4 pixel = read_imagef(source, sampler, sourceCoord);
102
103	pixel = (float4)(clLookUpTable[(int)(round(pixel.x * 255))],
104	clLookUpTable[(int)(round(pixel.y * 255))],
105	clLookUpTable[(int)(round(pixel.z * 255))],
106	pixel.w);
107
108	write_imagef(destination, sourceCoord, pixel);
109	}
110	);
111
112	cl_mem FilterContextOpenCL::OpenCLTransformColorSpace(cl_mem source, ColorSpace resultColorSpace, ColorSpace dstColorSpace, IntRect sourceSize)
113	{
114	DEFINE_STATIC_LOCAL(cl_mem, deviceRgbLUT, ());
115	DEFINE_STATIC_LOCAL(cl_mem, linearRgbLUT, ());
116
117	if ((resultColorSpace != ColorSpaceLinearRGB && resultColorSpace != ColorSpaceDeviceRGB)
118	\|\| (dstColorSpace != ColorSpaceLinearRGB && dstColorSpace != ColorSpaceDeviceRGB))
119	return source;
120
121
122	FilterContextOpenCL* context = FilterContextOpenCL::context();
123
124	cl_image_format clImageFormat;
125	clImageFormat.image_channel_order = CL_RGBA;
126	clImageFormat.image_channel_data_type = CL_UNORM_INT8;
127
128	cl_mem destination = clCreateImage2D(context->deviceContext(), CL_MEM_READ_WRITE,
129	&clImageFormat, sourceSize.width(), sourceSize.height(), 0, 0, 0);
130
131	if(!m_OpenCLTransformColorSpaceProgram){
132	m_OpenCLTransformColorSpaceProgram = compileProgram(OpenCLTransformColorSpaceKernelProgram);
133	m_OpenCLTransformColorSpace = kernelByName(m_OpenCLTransformColorSpaceProgram, "OpenCLTransformColorSpace");
134	}
135
136	RunKernel kernel(context, m_OpenCLTransformColorSpace, sourceSize.width(), sourceSize.height());
137	kernel.addArgument(source);
138	kernel.addArgument(destination);
139
140	if (dstColorSpace == ColorSpaceLinearRGB) {
141	if (!linearRgbLUT) {
142	Vector<float> lookUpTable;
143	for (unsigned i = 0; i < 256; i++) {
144	float color = i / 255.0f;
145	color = (color <= 0.04045f ? color / 12.92f : pow((color + 0.055f) / 1.055f, 2.4f));
146	color = std::max(0.0f, color);
147	color = std::min(1.0f, color);
148	lookUpTable.append((round(color * 255)) / 255);
149	}
150	linearRgbLUT = context->uploadBuffer(context, lookUpTable.data(), sizeof(float) * 256);
151	}
152	kernel.addArgument(linearRgbLUT);
153	} else if (dstColorSpace == ColorSpaceDeviceRGB) {
154	if (!deviceRgbLUT) {
155	Vector<float> lookUpTable;
156	for (unsigned i = 0; i < 256; i++) {
157	float color = i / 255.0f;
158	color = (powf(color, 1.0f / 2.4f) * 1.055f) - 0.055f;
159	color = std::max(0.0f, color);
160	color = std::min(1.0f, color);
161	lookUpTable.append((round(color * 255)) / 255);
162	}
163	deviceRgbLUT = context->uploadBuffer(context, lookUpTable.data(), sizeof(float) * 256);
164	}
165	kernel.addArgument(deviceRgbLUT);
166	}
167
168	kernel.run();
169
170	return destination;
171	}
172
173	static const char* FillKernelProgram =
174	PROGRAM_STR(
175	__kernel void Fill(__write_only image2d_t destination,
176	float r,
177	float g,
178	float b,
179	float a)
180	{
181	float4 sourcePixel = (float4)(r,g,b,a);
182	write_imagef(destination, (int2)(get_global_id(0), get_global_id(1)), sourcePixel);
183	}
184	);
185
186	void FilterContextOpenCL::Fill(cl_mem image, IntSize imageSize, Color color)
187	{
188
189	if (!m_Fill) {
190	m_FillProgram = compileProgram(FillKernelProgram);
191	ASSERT(m_FillProgram);
192	m_Fill = kernelByName(m_FillProgram, "Fill");
193	ASSERT(m_Fill);
194	}
195
196	float r,g,b,a;
197
198	color.getRGBA(r,g,b,a);
199
200	RunKernel kernel(this, m_Fill, imageSize.width(), imageSize.height());
201	kernel.addArgument(image);
202	kernel.addArgument(r);
203	kernel.addArgument(g);
204	kernel.addArgument(b);
205	kernel.addArgument(a);
206	kernel.run();
207
208	}
209
210	cl_mem FilterContextOpenCL::uploadBuffer(WebCore::FilterContextOpenCL* context, void* buffer, int size)
211	{
212	cl_int err;
213	cl_mem result = clCreateBuffer(context->deviceContext(), CL_MEM_READ_ONLY, size, 0, &err);
214	if (err != CL_SUCCESS)
215	OpenCLPrintError(err);
216
217	err = clEnqueueWriteBuffer(context->commandQueue(), result, CL_TRUE, 0, size, buffer, 0, 0, 0);
218	if (err != CL_SUCCESS)
219	OpenCLPrintError(err);
220
221	return result;
222	}
223
224	cl_mem FilterContextOpenCL::createOpenCLImage(FloatSize paintSize)
225	{
226	FilterContextOpenCL* context = FilterContextOpenCL::context();
227
228	cl_image_format clImageFormat;
229	clImageFormat.image_channel_order = CL_RGBA;
230	clImageFormat.image_channel_data_type = CL_UNORM_INT8;
231
232	cl_int er;
233	cl_mem image = clCreateImage2D(context->deviceContext(), CL_MEM_READ_WRITE ,
234	&clImageFormat, paintSize.width(), paintSize.height(), 0, 0, &er);
235	if (er != CL_SUCCESS)
236	context->OpenCLPrintError(er);
237	context->OpenCLDebugImage(image);
238	return image;
239	}
240
241	void FilterContextOpenCL::OpenCLPrintError(cl_int error)
242	{
243	switch (error) {
244	case CL_SUCCESS:
245	fprintf(stderr, "CL_SUCCESS\n");
246	break;
247	case CL_DEVICE_NOT_FOUND:
248	fprintf(stderr, "CL_DEVICE_NOT_FOUND\n");
249	break;
250	case CL_DEVICE_NOT_AVAILABLE:
251	fprintf(stderr, "CL_DEVICE_NOT_AVAILABLE\n");
252	break;
253	case CL_COMPILER_NOT_AVAILABLE:
254	fprintf(stderr, "CL_COMPILER_NOT_AVAILABLE\n");
255	break;
256	case CL_MEM_OBJECT_ALLOCATION_FAILURE:
257	fprintf(stderr, "CL_MEM_OBJECT_ALLOCATION_FAILURE\n");
258	break;
259	case CL_OUT_OF_RESOURCES:
260	fprintf(stderr, "CL_OUT_OF_RESOURCES\n");
261	break;
262	case CL_OUT_OF_HOST_MEMORY:
263	fprintf(stderr, "CL_OUT_OF_HOST_MEMORY\n");
264	break;
265	case CL_PROFILING_INFO_NOT_AVAILABLE:
266	fprintf(stderr, "CL_PROFILING_INFO_NOT_AVAILABLE\n");
267	break;
268	case CL_MEM_COPY_OVERLAP:
269	fprintf(stderr, "CL_MEM_COPY_OVERLAP\n");
270	break;
271	case CL_IMAGE_FORMAT_MISMATCH:
272	fprintf(stderr, "CL_IMAGE_FORMAT_MISMATCH\n");
273	break;
274	case CL_IMAGE_FORMAT_NOT_SUPPORTED:
275	fprintf(stderr, "CL_IMAGE_FORMAT_NOT_SUPPORTED\n");
276	break;
277	case CL_BUILD_PROGRAM_FAILURE:
278	fprintf(stderr, "CL_BUILD_PROGRAM_FAILURE\n");
279	break;
280	case CL_MAP_FAILURE:
281	fprintf(stderr, "CL_MAP_FAILURE\n");
282	break;
283	case CL_INVALID_VALUE:
284	fprintf(stderr, "CL_INVALID_VALUE\n");
285	break;
286	case CL_INVALID_DEVICE_TYPE:
287	fprintf(stderr, "CL_INVALID_DEVICE_TYPE\n");
288	break;
289	case CL_INVALID_PLATFORM:
290	fprintf(stderr, "CL_INVALID_PLATFORM\n");
291	break;
292	case CL_INVALID_DEVICE:
293	fprintf(stderr, "CL_INVALID_DEVICE\n");
294	break;
295	case CL_INVALID_CONTEXT:
296	fprintf(stderr, "CL_INVALID_CONTEXT\n");
297	break;
298	case CL_INVALID_QUEUE_PROPERTIES:
299	fprintf(stderr, "CL_INVALID_QUEUE_PROPERTIES\n");
300	break;
301	case CL_INVALID_COMMAND_QUEUE:
302	fprintf(stderr, "CL_INVALID_COMMAND_QUEUE\n");
303	break;
304	case CL_INVALID_HOST_PTR:
305	fprintf(stderr, "CL_INVALID_HOST_PTR\n");
306	break;
307	case CL_INVALID_MEM_OBJECT:
308	fprintf(stderr, "CL_INVALID_MEM_OBJECT\n");
309	break;
310	case CL_INVALID_IMAGE_FORMAT_DESCRIPTOR:
311	fprintf(stderr, "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR\n");
312	break;
313	case CL_INVALID_IMAGE_SIZE:
314	fprintf(stderr, "CL_INVALID_IMAGE_SIZE\n");
315	break;
316	case CL_INVALID_SAMPLER:
317	fprintf(stderr, "CL_INVALID_SAMPLER\n");
318	break;
319	case CL_INVALID_BINARY:
320	fprintf(stderr, "CL_INVALID_BINARY\n");
321	break;
322	case CL_INVALID_BUILD_OPTIONS:
323	fprintf(stderr, "CL_INVALID_BUILD_OPTIONS\n");
324	break;
325	case CL_INVALID_PROGRAM:
326	fprintf(stderr, "CL_INVALID_PROGRAM\n");
327	break;
328	case CL_INVALID_PROGRAM_EXECUTABLE:
329	fprintf(stderr, "CL_INVALID_PROGRAM_EXECUTABLE\n");
330	break;
331	case CL_INVALID_KERNEL_NAME:
332	fprintf(stderr, "CL_INVALID_KERNEL_NAME\n");
333	break;
334	case CL_INVALID_KERNEL_DEFINITION:
335	fprintf(stderr, "CL_INVALID_KERNEL_DEFINITION\n");
336	break;
337	case CL_INVALID_KERNEL:
338	fprintf(stderr, "CL_INVALID_KERNEL\n");
339	break;
340	case CL_INVALID_ARG_INDEX:
341	fprintf(stderr, "CL_INVALID_ARG_INDEX\n");
342	break;
343	case CL_INVALID_ARG_VALUE:
344	fprintf(stderr, "CL_INVALID_ARG_VALUE\n");
345	break;
346	case CL_INVALID_ARG_SIZE:
347	fprintf(stderr, "CL_INVALID_ARG_SIZE\n");
348	break;
349	case CL_INVALID_KERNEL_ARGS:
350	fprintf(stderr, "CL_INVALID_KERNEL_ARGS\n");
351	break;
352	case CL_INVALID_WORK_DIMENSION:
353	fprintf(stderr, "CL_INVALID_WORK_DIMENSION\n");
354	break;
355	case CL_INVALID_WORK_GROUP_SIZE:
356	fprintf(stderr, "CL_INVALID_WORK_GROUP_SIZE\n");
357	break;
358	case CL_INVALID_WORK_ITEM_SIZE:
359	fprintf(stderr, "CL_INVALID_WORK_ITEM_SIZE\n");
360	break;
361	case CL_INVALID_GLOBAL_OFFSET:
362	fprintf(stderr, "CL_INVALID_GLOBAL_OFFSET\n");
363	break;
364	case CL_INVALID_EVENT_WAIT_LIST:
365	fprintf(stderr, "CL_INVALID_EVENT_WAIT_LIST\n");
366	break;
367	case CL_INVALID_EVENT:
368	fprintf(stderr, "CL_INVALID_EVENT\n");
369	break;
370	case CL_INVALID_OPERATION:
371	fprintf(stderr, "CL_INVALID_OPERATION\n");
372	break;
373	case CL_INVALID_GL_OBJECT:
374	fprintf(stderr, "CL_INVALID_GL_OBJECT\n");
375	break;
376	case CL_INVALID_BUFFER_SIZE:
377	fprintf(stderr, "CL_INVALID_BUFFER_SIZE\n");
378	break;
379	case CL_INVALID_MIP_LEVEL:
380	fprintf(stderr, "CL_INVALID_MIP_LEVEL\n");
381	break;
382	case CL_INVALID_GLOBAL_WORK_SIZE:
383	fprintf(stderr, "CL_INVALID_GLOBAL_WORK_SIZE\n");
384	break;
385	default:
386	fprintf(stderr, "Unknown error code : %d\n", error);
387	break;
388	}
389	}
390
391	void FilterContextOpenCL::OpenCLDebugImage(cl_mem image)
392	{
393	size_t width;
394	size_t height;
395	cl_int err;
396	size_t channelSize;
397
398	err = clGetImageInfo(image, CL_IMAGE_WIDTH, sizeof(width), &width, NULL);
399	if (err != CL_SUCCESS)
400	OpenCLPrintError(err);
401
402	err = clGetImageInfo(image, CL_IMAGE_HEIGHT, sizeof(height), &height, NULL);
403	if (err != CL_SUCCESS)
404	OpenCLPrintError(err);
405
406	err = clGetImageInfo(image, CL_IMAGE_ELEMENT_SIZE, sizeof(channelSize), &channelSize, NULL);
407	if (err != CL_SUCCESS)
408	OpenCLPrintError(err);
409	if (err != CL_SUCCESS)
410	printf("OpenCLDebugImage: width: %ld height: %ld channelSize: %ld \n", width, height, channelSize);
411	}
412
413
414
415
416	} // namespace WebCore
417
418	#endif

Line 0 a/Source/WebCore/platform/graphics/filters/OpenCL/OpenCLContext.h_sec1
- a/Source/WebCore/platform/graphics/filters/OpenCL/OpenCLContext.h +193 lines
		1	#ifndef OpenCLContext_h
		2	#define OpenCLContext_h
		3
		4	#include "config.h"
		5	#if ENABLE(OPENCL)
		6
		7	#include "CL/cl.h"
		8	#include "Color.h"
		9	#include "ColorSpace.h"
		10	#include "IntSize.h"
		11	#include "IntRect.h"
		12	#include "FloatPoint.h"
		13	#include "FloatSize.h"
		14	#include "stdio.h"
		15
		16	namespace WebCore {
		17
		18	class FilterContextOpenCL
		19	{
		20	public:
		21	// Returns NULL if initialization failed.
		22	static FilterContextOpenCL* context();
		23
		24	cl_context deviceContext() { return m_deviceContext; }
		25	cl_command_queue commandQueue() { return m_commandQueue; }
		26
		27	cl_mem createOpenCLImage(FloatSize);
		28
		29	static cl_mem uploadBuffer(WebCore::FilterContextOpenCL* context, void* buffer, int size);
		30	cl_mem OpenCLTransformColorSpace(cl_mem, ColorSpace, ColorSpace, IntRect);
		31	void Fill(cl_mem, IntSize, Color);
		32
		33	inline void compileFEComposite();
		34	inline void compileFEOffset();
		35	inline void compileFEBlend();
		36	inline void compileFEGaussianBlur();
		37	inline void compileFETurbulence();
		38	inline void compileFEColorMatrix();
		39	inline void compileFEConvolveMatrix();
		40	inline void compileFEComponentTransfer();
		41	inline void compileFEMerge();
		42	inline void compileFETile();
		43	inline void compileFEMorphology();
		44
		45	inline void applyFEColorMatrix(cl_mem, IntSize, cl_mem, FloatPoint, cl_mem, int);
		46	inline void applyFEColorMatrixLuminanceToAlpha(cl_mem, IntSize, cl_mem, FloatPoint, int);
		47	inline void applyFEConvolveMatrix(cl_mem, IntSize, cl_mem, IntSize, cl_mem, int, float, float, IntPoint, bool);
		48	inline void applyFETurbulence(cl_mem, IntSize, cl_mem, cl_mem, cl_mem, cl_mem, cl_mem, cl_mem, int, int, int, int, float, float, bool, int, int);
		49	inline void applyFEComposite(cl_mem, IntSize, cl_mem, cl_mem, FloatPoint, FloatPoint, float, float, float, float, int);
		50	inline void applyFEComponentTransfer(cl_mem, IntSize, cl_mem, cl_mem, FloatPoint);
		51	inline void applyFEOffset(cl_mem, IntSize, cl_mem, FloatPoint);
		52	inline void applyFEBlend(cl_mem, IntSize, cl_mem, cl_mem, int,FloatPoint, FloatPoint);
		53	inline void applyFEGaussianBlur(cl_mem, IntSize, cl_mem);
		54	inline void applyFEMerge(cl_mem, IntSize, cl_mem, FloatRect);
		55	inline void applyFETiling(cl_mem, IntSize, cl_mem, FloatRect, FloatPoint);
		56	inline void applyFETileBuilder(cl_mem, FloatRect, cl_mem, FloatSize);
		57	inline void applyFEMorphology(cl_mem, IntSize, cl_mem, int, int, int, int, int);
		58
		59	static void OpenCLPrintError(cl_int);
		60	static void OpenCLDebugImage(cl_mem);
		61
		62	private:
		63
		64	class RunKernel {
65	public:
66	RunKernel(FilterContextOpenCL* context, cl_kernel kernel, size_t width, size_t height)
67	: m_context(context)
68	, m_kernel(kernel)
69	, index(0)
70	{
71	m_globalSize[0] = width;
72	m_globalSize[1] = height;
73	}
74
75	void addArgument(cl_mem handle)
76	{
77	cl_int err;
78	err = clSetKernelArg(m_kernel, index++, sizeof(cl_mem), reinterpret_cast<void*>(&handle));
79	if (err != CL_SUCCESS){
80	m_context->OpenCLPrintError(err);
81	}
82	}
83
84	void addArgument(cl_int value)
85	{
86	cl_int err;
87	err = clSetKernelArg(m_kernel, index++, sizeof(cl_int), reinterpret_cast<void*>(&value));
88	if (err != CL_SUCCESS){
89	m_context->OpenCLPrintError(err);
90	}
91	}
92
93	void addArgument(cl_float value)
94	{
95	cl_int err;
96	err = clSetKernelArg(m_kernel, index++, sizeof(cl_float), reinterpret_cast<void*>(&value));
97	if (err != CL_SUCCESS){
98	m_context->OpenCLPrintError(err);
99	}
100	}
101
102	void addArgument(cl_sampler handle)
103	{
104	cl_int err;
105	err = clSetKernelArg(m_kernel, index++, sizeof(cl_sampler), reinterpret_cast<void*>(&handle));
106	if (err != CL_SUCCESS){
107	m_context->OpenCLPrintError(err);
108	}
109	}
110
111	void run()
112	{
113	clFinish(m_context->m_commandQueue);
114	clEnqueueNDRangeKernel(m_context->m_commandQueue, m_kernel, 2, 0, m_globalSize, 0, 0, 0, 0);
115	}
116
117	FilterContextOpenCL* m_context;
118	cl_kernel m_kernel;
119	size_t m_globalSize[2];
120	int index;
121	};
122
123	static cl_program compileProgram(const char*);
124	static inline cl_kernel kernelByName(cl_program program, const char* name) { return clCreateKernel(program, name, NULL); }
125
126	cl_command_queue m_commandQueue;
127	cl_device_id m_device;
128	cl_context m_deviceContext;
129
130	cl_program m_OpenCLTransformColorSpaceProgram;
131	cl_kernel m_OpenCLTransformColorSpace;
132
133	cl_program m_FEColorMatrixProgram;
134	cl_kernel m_FEMatrix;
135	cl_kernel m_FESaturate;
136	cl_kernel m_FEHuerotate;
137	cl_kernel m_FELuminance;
138
139	cl_program m_FEConvolveMatrixProgram;
140	cl_kernel m_FEConvolveMatrix;
141	cl_kernel m_FEConvolveMatrixAlpha;
142
143	cl_program m_FETurbulenceProgram;
144	cl_kernel m_FETurbulence;
145
146	cl_program m_feOffsetProgram;
147	cl_kernel m_feOffset;
148
149	cl_program m_feBlendProgram;
150	cl_kernel m_feNormalBlend;
151	cl_kernel m_feMultiplyBlend;
152	cl_kernel m_feScreenBlend;
153	cl_kernel m_feDarkenBlend;
154	cl_kernel m_feLightenBlend;
155	cl_kernel m_feUnknowBlend;
156
157	cl_program m_feGaussianBlurProgram;
158	cl_kernel m_feGaussianBlur;
159
160	cl_program m_FEFloodProgram;
161	cl_kernel m_FEFlood;
162
163	cl_program m_feCompositeProgram;
164	cl_kernel m_feArithmetic;
165	cl_kernel m_feAtop;
166	cl_kernel m_feIn;
167	cl_kernel m_feOut;
168	cl_kernel m_feOver;
169	cl_kernel m_feXor;
170
171	cl_program m_FEMergeProgram;
172	cl_kernel m_FEMerge;
173
174	cl_program m_feTileProgram;
175	cl_kernel m_feTile;
176	cl_kernel m_TileBuilder;
177
178	cl_program m_feMorphologyProgram;
179	cl_kernel m_feMorphologyDilate;
180	cl_kernel m_feMorphologyErode;
181
182	cl_program m_FillProgram;
183	cl_kernel m_Fill;
184
185	cl_program m_feComponentTransferProgram;
186	cl_kernel m_feComponentTransfer;
187	};
188
189	} // namespace WebCore
190
191	#endif // ENABLE(OPENCL)
192
193	#endif

Line 0 a/Source/WebCore/platform/graphics/filters/OpenCL/OpenCLFEBlend.cpp_sec1
- a/Source/WebCore/platform/graphics/filters/OpenCL/OpenCLFEBlend.cpp +240 lines
		1	/*
		2	* Copyright (C) 2012 Tamas Czene <tczene@inf.u-szeged.hu>
		3	*
		4	* This library is free software; you can redistribute it and/or
		5	* modify it under the terms of the GNU Library General Public
		6	* License as published by the Free Software Foundation; either
		7	* version 2 of the License, or (at your option) any later version.
		8	*
		9	* This library is distributed in the hope that it will be useful,
		10	* but WITHOUT ANY WARRANTY; without even the implied warranty of
		11	* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
		12	* Library General Public License for more details.
		13	*
		14	* You should have received a copy of the GNU Library General Public License
		15	* along with this library; see the file COPYING.LIB. If not, write to
		16	* the Free Software Foundation, Inc., 51 Franklin Street, Fifth Floor,
		17	* Boston, MA 02110-1301, USA.
		18	*/
		19
		20	#include "config.h"
		21	#if ENABLE(FILTERS) && ENABLE(OPENCL)
		22	#include "FEBlend.h"
		23	#include "OpenCLContext.h"
		24
		25	#include <stdio.h>
		26
		27	#define PROGRAM_STR(Src) #Src
		28	#define PROGRAM(Src) PROGRAM_STR(Src)
		29
		30	namespace WebCore {
		31
		32	static const char* FEBlendKernelProgram =
		33	PROGRAM_STR(
		34	const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE \| CLK_ADDRESS_CLAMP \| CLK_FILTER_NEAREST;
		35
		36	__kernel void NormalBlend(__write_only image2d_t destination,
		37	__read_only image2d_t sourceA,
		38	__read_only image2d_t sourceB,
		39	float xA,
		40	float yA,
		41	float xB,
		42	float yB)
		43	{
		44	int2 sourceCoordA = (int2) (get_global_id(0) + xA, get_global_id(1) + yA);
		45	int2 sourceCoordB = (int2) (get_global_id(0) + xB, get_global_id(1) + yB);
		46	float4 sourcePixelA = read_imagef(sourceA, sampler, sourceCoordA);
		47	float4 sourcePixelB = read_imagef(sourceB, sampler, sourceCoordB);
		48	float4 destinationPixel = (float4)((1.0 - sourcePixelA.w) * sourcePixelB.x + sourcePixelA.x,
		49	(1.0 - sourcePixelA.w) * sourcePixelB.y + sourcePixelA.y,
		50	(1.0 - sourcePixelA.w) * sourcePixelB.z + sourcePixelA.z,
		51	1.0 - (1.0 - sourcePixelA.w) * (1.0 - sourcePixelB.w));
		52	write_imagef(destination, (int2)(get_global_id(0), get_global_id(1)), destinationPixel);
		53	}
		54
		55	__kernel void MultiplyBlend(__write_only image2d_t destination,
		56	__read_only image2d_t sourceA,
		57	__read_only image2d_t sourceB,
		58	float xA,
		59	float yA,
		60	float xB,
		61	float yB)
		62	{
		63	int2 sourceCoordA = (int2) (get_global_id(0) + xA, get_global_id(1) + yA);
		64	int2 sourceCoordB = (int2) (get_global_id(0) + xB, get_global_id(1) + yB);
65	float4 sourcePixelA = read_imagef(sourceA, sampler, sourceCoordA);
66	float4 sourcePixelB = read_imagef(sourceB, sampler, sourceCoordB);
67	float4 destinationPixel = (float4)((1.0 - sourcePixelA.w) * sourcePixelB.x + (1.0 - sourcePixelB.w) * sourcePixelA.x + sourcePixelA.x * sourcePixelB.x,
68	(1.0 - sourcePixelA.w) * sourcePixelB.y + (1.0 - sourcePixelB.w) * sourcePixelA.y + sourcePixelA.y * sourcePixelB.y,
69	(1.0 - sourcePixelA.w) * sourcePixelB.z + (1.0 - sourcePixelB.w) * sourcePixelA.z + sourcePixelA.z * sourcePixelB.z,
70	1.0 - (1.0 - sourcePixelA.w) * (1.0 - sourcePixelB.w));
71	write_imagef(destination, (int2)(get_global_id(0), get_global_id(1)), destinationPixel);
72	}
73
74	__kernel void ScreenBlend(__write_only image2d_t destination,
75	__read_only image2d_t sourceA,
76	__read_only image2d_t sourceB,
77	float xA,
78	float yA,
79	float xB,
80	float yB)
81	{
82	int2 sourceCoordA = (int2) (get_global_id(0) + xA, get_global_id(1) + yA);
83	int2 sourceCoordB = (int2) (get_global_id(0) + xB, get_global_id(1) + yB);
84	float4 sourcePixelA = read_imagef(sourceA, sampler, sourceCoordA);
85	float4 sourcePixelB = read_imagef(sourceB, sampler, sourceCoordB);
86	float4 destinationPixel = (float4)(sourcePixelB.x + sourcePixelA.x - sourcePixelA.x * sourcePixelB.x,
87	sourcePixelB.y + sourcePixelA.y - sourcePixelA.y * sourcePixelB.y,
88	sourcePixelB.z + sourcePixelA.z - sourcePixelA.z * sourcePixelB.z,
89	1.0 - (1.0 - sourcePixelA.w) * (1.0 - sourcePixelB.w));
90	write_imagef(destination, (int2)(get_global_id(0), get_global_id(1)), destinationPixel);
91	}
92
93	__kernel void DarkenBlend(__write_only image2d_t destination,
94	__read_only image2d_t sourceA,
95	__read_only image2d_t sourceB,
96	float xA,
97	float yA,
98	float xB,
99	float yB)
100	{
101	int2 sourceCoordA = (int2) (get_global_id(0) + xA, get_global_id(1) + yA);
102	int2 sourceCoordB = (int2) (get_global_id(0) + xB, get_global_id(1) + yB);
103	float4 sourcePixelA = read_imagef(sourceA, sampler, sourceCoordA);
104	float4 sourcePixelB = read_imagef(sourceB, sampler, sourceCoordB);
105	float4 destinationPixel = (float4)(min((1.0 - sourcePixelA.w) * sourcePixelB.x + sourcePixelA.x, (1.0 - sourcePixelB.w) * sourcePixelA.x + sourcePixelB.x),
106	min((1.0 - sourcePixelA.w) * sourcePixelB.y + sourcePixelA.y, (1.0 - sourcePixelB.w) * sourcePixelA.y + sourcePixelB.y),
107	min((1.0 - sourcePixelA.w) * sourcePixelB.z + sourcePixelA.z, (1.0 - sourcePixelB.w) * sourcePixelA.z + sourcePixelB.z),
108	1.0 - (1.0 - sourcePixelA.w) * (1.0 - sourcePixelB.w));
109	write_imagef(destination, (int2)(get_global_id(0), get_global_id(1)), destinationPixel);
110	}
111
112	__kernel void LightenBlend(__write_only image2d_t destination,
113	__read_only image2d_t sourceA,
114	__read_only image2d_t sourceB,
115	float xA,
116	float yA,
117	float xB,
118	float yB)
119	{
120	int2 sourceCoordA = (int2) (get_global_id(0) + xA, get_global_id(1) + yA);
121	int2 sourceCoordB = (int2) (get_global_id(0) + xB, get_global_id(1) + yB);
122	float4 sourcePixelA = read_imagef(sourceA, sampler, sourceCoordA);
123	float4 sourcePixelB = read_imagef(sourceB, sampler, sourceCoordB);
124	float4 destinationPixel = (float4)(max((1.0 - sourcePixelA.w) * sourcePixelB.x + sourcePixelA.x, (1.0 - sourcePixelB.w) * sourcePixelA.x + sourcePixelB.x),
125	max((1.0 - sourcePixelA.w) * sourcePixelB.y + sourcePixelA.y, (1.0 - sourcePixelB.w) * sourcePixelA.y + sourcePixelB.y),
126	max((1.0 - sourcePixelA.w) * sourcePixelB.z + sourcePixelA.z, (1.0 - sourcePixelB.w) * sourcePixelA.z + sourcePixelB.z),
127	1.0 - (1.0 - sourcePixelA.w) * (1.0 - sourcePixelB.w));
128	write_imagef(destination, (int2)(get_global_id(0), get_global_id(1)), destinationPixel);
129	}
130
131	__kernel void UnknowBlend(__write_only image2d_t destination,
132	__read_only image2d_t sourceA,
133	__read_only image2d_t sourceB,
134	float xA,
135	float yA,
136	float xB,
137	float yB)
138	{
139	int2 sourceCoord = (int2) (get_global_id(0), get_global_id(1));
140	float4 destinationPixel = (float4)(0.0,0.0,0.0,0.0);
141	write_imagef(destination, (int2)(get_global_id(0), get_global_id(1)), destinationPixel);
142	}
143	);
144
145	inline void FilterContextOpenCL::compileFEBlend()
146	{
147	if (!m_feNormalBlend) {
148	m_feBlendProgram = compileProgram(FEBlendKernelProgram);
149	ASSERT(m_feBlendProgram);
150
151	m_feNormalBlend = kernelByName(m_feBlendProgram, "NormalBlend");
152	ASSERT(m_feNormalBlend);
153	m_feMultiplyBlend = kernelByName(m_feBlendProgram, "MultiplyBlend");
154	ASSERT(m_feNormalBlend);
155	m_feScreenBlend = kernelByName(m_feBlendProgram, "ScreenBlend");
156	ASSERT(m_feNormalBlend);
157	m_feDarkenBlend = kernelByName(m_feBlendProgram, "DarkenBlend");
158	ASSERT(m_feNormalBlend);
159	m_feLightenBlend = kernelByName(m_feBlendProgram, "LightenBlend");
160	ASSERT(m_feNormalBlend);
161	m_feUnknowBlend = kernelByName(m_feBlendProgram, "UnknowBlend");
162	ASSERT(m_feNormalBlend);
163	}
164	}
165
166	inline void FilterContextOpenCL::applyFEBlend(cl_mem destination, IntSize destinationSize, cl_mem sourceA, cl_mem sourceB, int mode, FloatPoint relativeSourceLocationA, FloatPoint relativeSourceLocationB)
167	{
168	cl_kernel blendMode;
169
170	switch (mode) {
171	case FEBLEND_MODE_NORMAL:
172	blendMode = m_feNormalBlend;
173	break;
174	case FEBLEND_MODE_MULTIPLY:
175	blendMode = m_feMultiplyBlend;
176	break;
177	case FEBLEND_MODE_SCREEN:
178	blendMode = m_feScreenBlend;
179	break;
180	case FEBLEND_MODE_DARKEN:
181	blendMode = m_feDarkenBlend;
182	break;
183	case FEBLEND_MODE_LIGHTEN:
184	blendMode = m_feLightenBlend;
185	break;
186	case FEBLEND_MODE_UNKNOWN:
187	default:
188	blendMode = m_feUnknowBlend;
189	break;
190	}
191
192	float xA = relativeSourceLocationA.x();
193	float yA = relativeSourceLocationA.y();
194
195	float xB = relativeSourceLocationB.x();
196	float yB = relativeSourceLocationB.y();
197
198	RunKernel kernel(this, blendMode, destinationSize.width(), destinationSize.height());
199	kernel.addArgument(destination);
200	kernel.addArgument(sourceA);
201	kernel.addArgument(sourceB);
202	kernel.addArgument(xA);
203	kernel.addArgument(yA);
204	kernel.addArgument(xB);
205	kernel.addArgument(yB);
206	kernel.run();
207	}
208
209	bool FEBlend::platformApplyOpenCL()
210	{
211	FilterContextOpenCL* context = FilterContextOpenCL::context();
212
213	if (!context)
214	return false;
215
216	context->compileFEBlend();
217
218	FilterEffect* in = inputEffect(0);
219	FilterEffect* in2 = inputEffect(1);
220
221	FloatPoint relativeSourceLocationA(in->absolutePaintRect().location());
222	relativeSourceLocationA.setX(absolutePaintRect().x() - relativeSourceLocationA.x());
223	relativeSourceLocationA.setY(absolutePaintRect().y() - relativeSourceLocationA.y());
224
225	FloatPoint relativeSourceLocationB(in2->absolutePaintRect().location());
226	relativeSourceLocationB.setX(absolutePaintRect().x() - relativeSourceLocationB.x());
227	relativeSourceLocationB.setY(absolutePaintRect().y() - relativeSourceLocationB.y());
228
229	cl_mem sourceA = in->asOpenCLImage();
230	cl_mem sourceB = in2->asOpenCLImage();
231	cl_mem destination = createOpenCLImageResult();
232
233	context->applyFEBlend(destination, absolutePaintRect().size(), sourceA, sourceB, m_mode, relativeSourceLocationA, relativeSourceLocationB);
234
235	return true;
236	}
237
238	} // namespace WebCore
239
240	#endif // ENABLE(FILTERS) && ENABLE(OPENCL)