Removing dead code in platform/graphics/gpu
authormmaxfield@apple.com <mmaxfield@apple.com@268f45cc-cd09-0410-ab3c-d52691b4dbfc>
Fri, 20 Jun 2014 21:21:00 +0000 (21:21 +0000)
committermmaxfield@apple.com <mmaxfield@apple.com@268f45cc-cd09-0410-ab3c-d52691b4dbfc>
Fri, 20 Jun 2014 21:21:00 +0000 (21:21 +0000)
https://bugs.webkit.org/show_bug.cgi?id=134065

Reviewed by Simon Fraser.

No new tests because there is no behavior change.

* WebCore.vcxproj/WebCore.vcxproj:
* WebCore.vcxproj/WebCore.vcxproj.filters:
* WebCore.xcodeproj/project.pbxproj:
* platform/graphics/gpu/LoopBlinnClassifier.cpp: Removed.
* platform/graphics/gpu/LoopBlinnClassifier.h: Removed.
* platform/graphics/gpu/LoopBlinnConstants.h: Removed.
* platform/graphics/gpu/LoopBlinnLocalTriangulator.cpp: Removed.
* platform/graphics/gpu/LoopBlinnLocalTriangulator.h: Removed.
* platform/graphics/gpu/LoopBlinnMathUtils.cpp: Removed.
* platform/graphics/gpu/LoopBlinnMathUtils.h: Removed.
* platform/graphics/gpu/LoopBlinnPathCache.cpp: Removed.
* platform/graphics/gpu/LoopBlinnPathCache.h: Removed.
* platform/graphics/gpu/LoopBlinnPathProcessor.h: Removed.
* platform/graphics/gpu/LoopBlinnShader.cpp: Removed.
* platform/graphics/gpu/LoopBlinnShader.h: Removed.
* platform/graphics/gpu/LoopBlinnSolidFillShader.cpp: Removed.
* platform/graphics/gpu/LoopBlinnSolidFillShader.h: Removed.
* platform/graphics/gpu/LoopBlinnTextureCoords.cpp: Removed.
* platform/graphics/gpu/LoopBlinnTextureCoords.h: Removed.
* platform/graphics/gpu/Shader.cpp: Removed.
* platform/graphics/gpu/Shader.h: Removed.
* platform/graphics/gpu/SharedGraphicsContext3D.cpp: Removed.
* platform/graphics/gpu/SharedGraphicsContext3D.h: Removed.
* platform/graphics/gpu/mac/DrawingBufferMac.mm: Removed.
* platform/graphics/gpu/opencl/FilterContextOpenCL.cpp: Removed.
* platform/graphics/gpu/opencl/FilterContextOpenCL.h: Removed.
* platform/graphics/gpu/opencl/OpenCLFEColorMatrix.cpp: Removed.
* platform/graphics/gpu/opencl/OpenCLFEFlood.cpp: Removed.
* platform/graphics/gpu/opencl/OpenCLFEImage.cpp: Removed.
* platform/graphics/gpu/opencl/OpenCLFEMerge.cpp: Removed.
* platform/graphics/gpu/opencl/OpenCLFESourceAlpha.cpp: Removed.
* platform/graphics/gpu/opencl/OpenCLFESourceGraphic.cpp: Removed.
* platform/graphics/gpu/opencl/OpenCLFETurbulence.cpp: Removed.
* platform/graphics/gpu/opencl/OpenCLHandle.h: Removed.

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

34 files changed:
Source/WebCore/ChangeLog
Source/WebCore/WebCore.vcxproj/WebCore.vcxproj
Source/WebCore/WebCore.vcxproj/WebCore.vcxproj.filters
Source/WebCore/WebCore.xcodeproj/project.pbxproj
Source/WebCore/platform/graphics/gpu/LoopBlinnClassifier.cpp [deleted file]
Source/WebCore/platform/graphics/gpu/LoopBlinnClassifier.h [deleted file]
Source/WebCore/platform/graphics/gpu/LoopBlinnConstants.h [deleted file]
Source/WebCore/platform/graphics/gpu/LoopBlinnLocalTriangulator.cpp [deleted file]
Source/WebCore/platform/graphics/gpu/LoopBlinnLocalTriangulator.h [deleted file]
Source/WebCore/platform/graphics/gpu/LoopBlinnMathUtils.cpp [deleted file]
Source/WebCore/platform/graphics/gpu/LoopBlinnMathUtils.h [deleted file]
Source/WebCore/platform/graphics/gpu/LoopBlinnPathCache.cpp [deleted file]
Source/WebCore/platform/graphics/gpu/LoopBlinnPathCache.h [deleted file]
Source/WebCore/platform/graphics/gpu/LoopBlinnPathProcessor.h [deleted file]
Source/WebCore/platform/graphics/gpu/LoopBlinnShader.cpp [deleted file]
Source/WebCore/platform/graphics/gpu/LoopBlinnShader.h [deleted file]
Source/WebCore/platform/graphics/gpu/LoopBlinnSolidFillShader.cpp [deleted file]
Source/WebCore/platform/graphics/gpu/LoopBlinnSolidFillShader.h [deleted file]
Source/WebCore/platform/graphics/gpu/LoopBlinnTextureCoords.cpp [deleted file]
Source/WebCore/platform/graphics/gpu/LoopBlinnTextureCoords.h [deleted file]
Source/WebCore/platform/graphics/gpu/Shader.cpp [deleted file]
Source/WebCore/platform/graphics/gpu/Shader.h [deleted file]
Source/WebCore/platform/graphics/gpu/SharedGraphicsContext3D.cpp [deleted file]
Source/WebCore/platform/graphics/gpu/SharedGraphicsContext3D.h [deleted file]
Source/WebCore/platform/graphics/gpu/opencl/FilterContextOpenCL.cpp [deleted file]
Source/WebCore/platform/graphics/gpu/opencl/FilterContextOpenCL.h [deleted file]
Source/WebCore/platform/graphics/gpu/opencl/OpenCLFEColorMatrix.cpp [deleted file]
Source/WebCore/platform/graphics/gpu/opencl/OpenCLFEFlood.cpp [deleted file]
Source/WebCore/platform/graphics/gpu/opencl/OpenCLFEImage.cpp [deleted file]
Source/WebCore/platform/graphics/gpu/opencl/OpenCLFEMerge.cpp [deleted file]
Source/WebCore/platform/graphics/gpu/opencl/OpenCLFESourceAlpha.cpp [deleted file]
Source/WebCore/platform/graphics/gpu/opencl/OpenCLFESourceGraphic.cpp [deleted file]
Source/WebCore/platform/graphics/gpu/opencl/OpenCLFETurbulence.cpp [deleted file]
Source/WebCore/platform/graphics/gpu/opencl/OpenCLHandle.h [deleted file]

index 0e415cd404fde7349d7ec452aa065b972d74696c..ae940d62496b1db6696e0e1f118db14c3c4e137e 100644 (file)
@@ -1,3 +1,47 @@
+2014-06-20  Myles C. Maxfield  <mmaxfield@apple.com>
+
+        Removing dead code in platform/graphics/gpu
+        https://bugs.webkit.org/show_bug.cgi?id=134065
+
+        Reviewed by Simon Fraser.
+
+        No new tests because there is no behavior change.
+
+        * WebCore.vcxproj/WebCore.vcxproj:
+        * WebCore.vcxproj/WebCore.vcxproj.filters:
+        * WebCore.xcodeproj/project.pbxproj:
+        * platform/graphics/gpu/LoopBlinnClassifier.cpp: Removed.
+        * platform/graphics/gpu/LoopBlinnClassifier.h: Removed.
+        * platform/graphics/gpu/LoopBlinnConstants.h: Removed.
+        * platform/graphics/gpu/LoopBlinnLocalTriangulator.cpp: Removed.
+        * platform/graphics/gpu/LoopBlinnLocalTriangulator.h: Removed.
+        * platform/graphics/gpu/LoopBlinnMathUtils.cpp: Removed.
+        * platform/graphics/gpu/LoopBlinnMathUtils.h: Removed.
+        * platform/graphics/gpu/LoopBlinnPathCache.cpp: Removed.
+        * platform/graphics/gpu/LoopBlinnPathCache.h: Removed.
+        * platform/graphics/gpu/LoopBlinnPathProcessor.h: Removed.
+        * platform/graphics/gpu/LoopBlinnShader.cpp: Removed.
+        * platform/graphics/gpu/LoopBlinnShader.h: Removed.
+        * platform/graphics/gpu/LoopBlinnSolidFillShader.cpp: Removed.
+        * platform/graphics/gpu/LoopBlinnSolidFillShader.h: Removed.
+        * platform/graphics/gpu/LoopBlinnTextureCoords.cpp: Removed.
+        * platform/graphics/gpu/LoopBlinnTextureCoords.h: Removed.
+        * platform/graphics/gpu/Shader.cpp: Removed.
+        * platform/graphics/gpu/Shader.h: Removed.
+        * platform/graphics/gpu/SharedGraphicsContext3D.cpp: Removed.
+        * platform/graphics/gpu/SharedGraphicsContext3D.h: Removed.
+        * platform/graphics/gpu/mac/DrawingBufferMac.mm: Removed.
+        * platform/graphics/gpu/opencl/FilterContextOpenCL.cpp: Removed.
+        * platform/graphics/gpu/opencl/FilterContextOpenCL.h: Removed.
+        * platform/graphics/gpu/opencl/OpenCLFEColorMatrix.cpp: Removed.
+        * platform/graphics/gpu/opencl/OpenCLFEFlood.cpp: Removed.
+        * platform/graphics/gpu/opencl/OpenCLFEImage.cpp: Removed.
+        * platform/graphics/gpu/opencl/OpenCLFEMerge.cpp: Removed.
+        * platform/graphics/gpu/opencl/OpenCLFESourceAlpha.cpp: Removed.
+        * platform/graphics/gpu/opencl/OpenCLFESourceGraphic.cpp: Removed.
+        * platform/graphics/gpu/opencl/OpenCLFETurbulence.cpp: Removed.
+        * platform/graphics/gpu/opencl/OpenCLHandle.h: Removed.
+
 2014-06-20  Myles C. Maxfield  <mmaxfield@apple.com>
 
         [iOS] Support Khmer and Lao fallback fonts
index 1955aec4262a949f7cef3bfd6c2279b80ea0a740..46c4bce788c521fde84838c2743dcd8be8f3ecd6 100644 (file)
     </ClCompile>
     <ClCompile Include="..\platform\graphics\egl\GLContextEGL.cpp" />
     <ClCompile Include="..\platform\graphics\gpu\DrawingBuffer.cpp" />
-    <ClCompile Include="..\platform\graphics\gpu\LoopBlinnClassifier.cpp" />
-    <ClCompile Include="..\platform\graphics\gpu\LoopBlinnLocalTriangulator.cpp" />
-    <ClCompile Include="..\platform\graphics\gpu\LoopBlinnMathUtils.cpp" />
-    <ClCompile Include="..\platform\graphics\gpu\LoopBlinnPathCache.cpp" />
-    <ClCompile Include="..\platform\graphics\gpu\LoopBlinnShader.cpp" />
-    <ClCompile Include="..\platform\graphics\gpu\LoopBlinnSolidFillShader.cpp" />
-    <ClCompile Include="..\platform\graphics\gpu\LoopBlinnTextureCoords.cpp" />
-    <ClCompile Include="..\platform\graphics\gpu\Shader.cpp" />
-    <ClCompile Include="..\platform\graphics\gpu\SharedGraphicsContext3D.cpp" />
     <ClCompile Include="..\platform\graphics\gpu\Texture.cpp" />
     <ClCompile Include="..\platform\graphics\gpu\TilingData.cpp" />
     <ClCompile Include="..\platform\graphics\FormatConverter.cpp" />
     </ClInclude>
     <ClInclude Include="..\platform\graphics\egl\GLContextEGL.h" />
     <ClInclude Include="..\platform\graphics\gpu\DrawingBuffer.h" />
-    <ClInclude Include="..\platform\graphics\gpu\LoopBlinnClassifier.h" />
-    <ClInclude Include="..\platform\graphics\gpu\LoopBlinnConstants.h" />
-    <ClInclude Include="..\platform\graphics\gpu\LoopBlinnLocalTriangulator.h" />
-    <ClInclude Include="..\platform\graphics\gpu\LoopBlinnMathUtils.h" />
-    <ClInclude Include="..\platform\graphics\gpu\LoopBlinnPathCache.h" />
-    <ClInclude Include="..\platform\graphics\gpu\LoopBlinnPathProcessor.h" />
-    <ClInclude Include="..\platform\graphics\gpu\LoopBlinnShader.h" />
-    <ClInclude Include="..\platform\graphics\gpu\LoopBlinnSolidFillShader.h" />
-    <ClInclude Include="..\platform\graphics\gpu\LoopBlinnTextureCoords.h" />
-    <ClInclude Include="..\platform\graphics\gpu\Shader.h" />
-    <ClInclude Include="..\platform\graphics\gpu\SharedGraphicsContext3D.h" />
     <ClInclude Include="..\platform\graphics\gpu\Texture.h" />
     <ClInclude Include="..\platform\graphics\gpu\TilingData.h" />
     <ClInclude Include="..\platform\graphics\FormatConverter.h" />
   <ImportGroup Label="ExtensionTargets">
     <Import Project="$(VCTargetsPath)\BuildCustomizations\masm.targets" />
   </ImportGroup>
-</Project>
\ No newline at end of file
+</Project>
index 8a49a35d2c95660fe5c10697c502a7b86464d3a1..2bc62057f3d7697d0bb63774431902aa0ade7771 100644 (file)
     <ClCompile Include="..\platform\graphics\gpu\DrawingBuffer.cpp">
       <Filter>platform\graphics\gpu</Filter>
     </ClCompile>
-    <ClCompile Include="..\platform\graphics\gpu\LoopBlinnClassifier.cpp">
-      <Filter>platform\graphics\gpu</Filter>
-    </ClCompile>
-    <ClCompile Include="..\platform\graphics\gpu\LoopBlinnLocalTriangulator.cpp">
-      <Filter>platform\graphics\gpu</Filter>
-    </ClCompile>
-    <ClCompile Include="..\platform\graphics\gpu\LoopBlinnMathUtils.cpp">
-      <Filter>platform\graphics\gpu</Filter>
-    </ClCompile>
-    <ClCompile Include="..\platform\graphics\gpu\LoopBlinnPathCache.cpp">
-      <Filter>platform\graphics\gpu</Filter>
-    </ClCompile>
-    <ClCompile Include="..\platform\graphics\gpu\LoopBlinnShader.cpp">
-      <Filter>platform\graphics\gpu</Filter>
-    </ClCompile>
-    <ClCompile Include="..\platform\graphics\gpu\LoopBlinnSolidFillShader.cpp">
-      <Filter>platform\graphics\gpu</Filter>
-    </ClCompile>
-    <ClCompile Include="..\platform\graphics\gpu\LoopBlinnTextureCoords.cpp">
-      <Filter>platform\graphics\gpu</Filter>
-    </ClCompile>
-    <ClCompile Include="..\platform\graphics\gpu\Shader.cpp">
-      <Filter>platform\graphics\gpu</Filter>
-    </ClCompile>
-    <ClCompile Include="..\platform\graphics\gpu\SharedGraphicsContext3D.cpp">
-      <Filter>platform\graphics\gpu</Filter>
-    </ClCompile>
     <ClCompile Include="..\platform\graphics\gpu\Texture.cpp">
       <Filter>platform\graphics\gpu</Filter>
     </ClCompile>
     <ClInclude Include="..\platform\graphics\gpu\DrawingBuffer.h">
       <Filter>platform\graphics\gpu</Filter>
     </ClInclude>
-    <ClInclude Include="..\platform\graphics\gpu\LoopBlinnClassifier.h">
-      <Filter>platform\graphics\gpu</Filter>
-    </ClInclude>
-    <ClInclude Include="..\platform\graphics\gpu\LoopBlinnConstants.h">
-      <Filter>platform\graphics\gpu</Filter>
-    </ClInclude>
-    <ClInclude Include="..\platform\graphics\gpu\LoopBlinnLocalTriangulator.h">
-      <Filter>platform\graphics\gpu</Filter>
-    </ClInclude>
-    <ClInclude Include="..\platform\graphics\gpu\LoopBlinnMathUtils.h">
-      <Filter>platform\graphics\gpu</Filter>
-    </ClInclude>
-    <ClInclude Include="..\platform\graphics\gpu\LoopBlinnPathCache.h">
-      <Filter>platform\graphics\gpu</Filter>
-    </ClInclude>
-    <ClInclude Include="..\platform\graphics\gpu\LoopBlinnPathProcessor.h">
-      <Filter>platform\graphics\gpu</Filter>
-    </ClInclude>
-    <ClInclude Include="..\platform\graphics\gpu\LoopBlinnShader.h">
-      <Filter>platform\graphics\gpu</Filter>
-    </ClInclude>
-    <ClInclude Include="..\platform\graphics\gpu\LoopBlinnSolidFillShader.h">
-      <Filter>platform\graphics\gpu</Filter>
-    </ClInclude>
-    <ClInclude Include="..\platform\graphics\gpu\LoopBlinnTextureCoords.h">
-      <Filter>platform\graphics\gpu</Filter>
-    </ClInclude>
-    <ClInclude Include="..\platform\graphics\gpu\Shader.h">
-      <Filter>platform\graphics\gpu</Filter>
-    </ClInclude>
-    <ClInclude Include="..\platform\graphics\gpu\SharedGraphicsContext3D.h">
-      <Filter>platform\graphics\gpu</Filter>
-    </ClInclude>
     <ClInclude Include="..\platform\graphics\gpu\Texture.h">
       <Filter>platform\graphics\gpu</Filter>
     </ClInclude>
       <Filter>platform\win</Filter>
     </MASM>
   </ItemGroup>
-</Project>
\ No newline at end of file
+</Project>
index adc780fa358f1e5f349f9736aa0f7755e8e47e08..44b1d6bfa49e6f0e4131ed6b573d4ada0ae6f386 100644 (file)
                498391570F1E776900C23782 /* WebKitCSSMatrix.idl */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = text; path = WebKitCSSMatrix.idl; sourceTree = "<group>"; };
                498770C21242C535002226BA /* DrawingBuffer.cpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.cpp; name = DrawingBuffer.cpp; path = gpu/DrawingBuffer.cpp; sourceTree = "<group>"; };
                498770C31242C535002226BA /* DrawingBuffer.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; name = DrawingBuffer.h; path = gpu/DrawingBuffer.h; sourceTree = "<group>"; };
-               498770D01242C535002226BA /* Shader.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; name = Shader.h; path = gpu/Shader.h; sourceTree = "<group>"; };
                498770D71242C535002226BA /* Texture.cpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.cpp; name = Texture.cpp; path = gpu/Texture.cpp; sourceTree = "<group>"; };
                498770D81242C535002226BA /* Texture.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; name = Texture.h; path = gpu/Texture.h; sourceTree = "<group>"; };
                498770D91242C535002226BA /* TilingData.cpp */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.cpp.cpp; name = TilingData.cpp; path = gpu/TilingData.cpp; sourceTree = "<group>"; };
diff --git a/Source/WebCore/platform/graphics/gpu/LoopBlinnClassifier.cpp b/Source/WebCore/platform/graphics/gpu/LoopBlinnClassifier.cpp
deleted file mode 100644 (file)
index 672b4d7..0000000
+++ /dev/null
@@ -1,126 +0,0 @@
-/*
- * Copyright (C) 2010 Google Inc. All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions
- * are met:
- *
- * 1.  Redistributions of source code must retain the above copyright
- *     notice, this list of conditions and the following disclaimer.
- * 2.  Redistributions in binary form must reproduce the above copyright
- *     notice, this list of conditions and the following disclaimer in the
- *     documentation and/or other materials provided with the distribution.
- *
- * THIS SOFTWARE IS PROVIDED BY APPLE AND ITS CONTRIBUTORS "AS IS" AND ANY
- * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
- * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
- * DISCLAIMED. IN NO EVENT SHALL APPLE OR ITS CONTRIBUTORS BE LIABLE FOR ANY
- * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
- * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
- * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
- * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF
- * THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- */
-
-#include "config.h"
-
-#if ENABLE(ACCELERATED_2D_CANVAS)
-
-#include "LoopBlinnClassifier.h"
-
-#include "LoopBlinnMathUtils.h"
-
-namespace WebCore {
-
-using LoopBlinnMathUtils::approxEqual;
-using LoopBlinnMathUtils::roundToZero;
-
-LoopBlinnClassifier::Result LoopBlinnClassifier::classify(const FloatPoint& c0,
-                                                          const FloatPoint& c1,
-                                                          const FloatPoint& c2,
-                                                          const FloatPoint& c3)
-{
-    // Consult the chapter for the definitions of the following
-    // (terse) variable names. Note that the b0..b3 coordinates are
-    // homogeneous, so the "z" value (actually the w coordinate) must
-    // be 1.0.
-    FloatPoint3D b0(c0.x(), c0.y(), 1.0f);
-    FloatPoint3D b1(c1.x(), c1.y(), 1.0f);
-    FloatPoint3D b2(c2.x(), c2.y(), 1.0f);
-    FloatPoint3D b3(c3.x(), c3.y(), 1.0f);
-
-    // Compute a1..a3.
-    float a1 = b0 * b3.cross(b2);
-    float a2 = b1 * b0.cross(b3);
-    float a3 = b2 * b1.cross(b0);
-
-    // Compute d1..d3.
-    float d1 = a1 - 2 * a2 + 3 * a3;
-    float d2 = -a2 + 3 * a3;
-    float d3 = 3 * a3;
-
-    // Experimentation has shown that the texture coordinates computed
-    // from these values quickly become huge, leading to roundoff errors
-    // and artifacts in the shader. It turns out that if we normalize
-    // the vector defined by (d1, d2, d3), this fixes the problem of the
-    // texture coordinates getting too large without affecting the
-    // classification results.
-    FloatPoint3D nd(d1, d2, d3);
-    nd.normalize();
-    d1 = nd.x();
-    d2 = nd.y();
-    d3 = nd.z();
-
-    // Compute the discriminant.
-    // term0 is a common term in the computation which helps decide
-    // which way to classify the cusp case: as serpentine or loop.
-    float term0 = (3 * d2 * d2 - 4 * d1 * d3);
-    float discriminant = d1 * d1 * term0;
-
-    // Experimentation has also shown that when the classification is
-    // near the boundary between one curve type and another, the shader
-    // becomes numerically unstable, particularly with the cusp case.
-    // Correct for this by rounding d1..d3 and the discriminant to zero
-    // when they get near it.
-    d1 = roundToZero(d1);
-    d2 = roundToZero(d2);
-    d3 = roundToZero(d3);
-    discriminant = roundToZero(discriminant);
-
-    // Do the classification.
-    if (approxEqual(b0, b1) && approxEqual(b0, b2) && approxEqual(b0, b3))
-        return Result(kPoint, d1, d2, d3);
-
-    if (!discriminant) {
-        if (!d1 && !d2) {
-            if (!d3)
-                return Result(kLine, d1, d2, d3);
-            return Result(kQuadratic, d1, d2, d3);
-        }
-
-        if (!d1)
-            return Result(kCusp, d1, d2, d3);
-
-        // This is the boundary case described in Loop and Blinn's
-        // SIGGRAPH '05 paper of a cusp with inflection at infinity.
-        // Because term0 might not be exactly 0, we decide between using
-        // the serpentine and loop cases depending on its sign to avoid
-        // taking the square root of a negative number when computing the
-        // cubic texture coordinates.
-        if (term0 < 0)
-            return Result(kLoop, d1, d2, d3);
-
-        return Result(kSerpentine, d1, d2, d3);
-    }
-
-    if (discriminant > 0)
-        return Result(kSerpentine, d1, d2, d3);
-
-    // discriminant < 0
-    return Result(kLoop, d1, d2, d3);
-}
-
-} // namespace WebCore
-
-#endif
diff --git a/Source/WebCore/platform/graphics/gpu/LoopBlinnClassifier.h b/Source/WebCore/platform/graphics/gpu/LoopBlinnClassifier.h
deleted file mode 100644 (file)
index 8b7181b..0000000
+++ /dev/null
@@ -1,85 +0,0 @@
-/*
- * Copyright (C) 2010 Google Inc. All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions
- * are met:
- *
- * 1.  Redistributions of source code must retain the above copyright
- *     notice, this list of conditions and the following disclaimer.
- * 2.  Redistributions in binary form must reproduce the above copyright
- *     notice, this list of conditions and the following disclaimer in the
- *     documentation and/or other materials provided with the distribution.
- *
- * THIS SOFTWARE IS PROVIDED BY APPLE AND ITS CONTRIBUTORS "AS IS" AND ANY
- * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
- * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
- * DISCLAIMED. IN NO EVENT SHALL APPLE OR ITS CONTRIBUTORS BE LIABLE FOR ANY
- * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
- * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
- * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
- * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF
- * THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- */
-
-// Cubic curve classification algorithm from "Rendering Vector Art on
-// the GPU" by Loop and Blinn, GPU Gems 3, Chapter 25:
-// http://http.developer.nvidia.com/GPUGems3/gpugems3_ch25.html .
-
-#ifndef LoopBlinnClassifier_h
-#define LoopBlinnClassifier_h
-
-#include <wtf/Noncopyable.h>
-
-namespace WebCore {
-
-class FloatPoint;
-
-// Classifies cubic curves into specific types.
-class LoopBlinnClassifier {
-    WTF_MAKE_NONCOPYABLE(LoopBlinnClassifier);
-public:
-    // The types of cubic curves.
-    enum CurveType {
-        kSerpentine,
-        kCusp,
-        kLoop,
-        kQuadratic,
-        kLine,
-        kPoint
-    };
-
-    // The result of the classifier.
-    struct Result {
-    public:
-        Result(CurveType inputCurveType, float inputD1, float inputD2, float inputD3)
-            : curveType(inputCurveType)
-            , d1(inputD1)
-            , d2(inputD2)
-            , d3(inputD3) { }
-
-        CurveType curveType;
-
-        // These are coefficients used later in the computation of
-        // texture coordinates per vertex.
-        float d1;
-        float d2;
-        float d3;
-    };
-
-    // Classifies the given cubic bezier curve starting at c0, ending
-    // at c3, and affected by control points c1 and c2.
-    static Result classify(const FloatPoint& c0,
-                           const FloatPoint& c1,
-                           const FloatPoint& c2,
-                           const FloatPoint& c3);
-
-private:
-    // This class does not need to be instantiated.
-    LoopBlinnClassifier() { }
-};
-
-} // namespace WebCore
-
-#endif // LoopBlinnClassifier_h
diff --git a/Source/WebCore/platform/graphics/gpu/LoopBlinnConstants.h b/Source/WebCore/platform/graphics/gpu/LoopBlinnConstants.h
deleted file mode 100644 (file)
index 1997d9f..0000000
+++ /dev/null
@@ -1,40 +0,0 @@
-/*
- * Copyright (C) 2010 Google Inc. All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions
- * are met:
- *
- * 1.  Redistributions of source code must retain the above copyright
- *     notice, this list of conditions and the following disclaimer.
- * 2.  Redistributions in binary form must reproduce the above copyright
- *     notice, this list of conditions and the following disclaimer in the
- *     documentation and/or other materials provided with the distribution.
- *
- * THIS SOFTWARE IS PROVIDED BY APPLE AND ITS CONTRIBUTORS "AS IS" AND ANY
- * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
- * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
- * DISCLAIMED. IN NO EVENT SHALL APPLE OR ITS CONTRIBUTORS BE LIABLE FOR ANY
- * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
- * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
- * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
- * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF
- * THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- */
-
-#ifndef LoopBlinnConstants_h
-#define LoopBlinnConstants_h
-
-namespace WebCore {
-namespace LoopBlinnConstants {
-
-enum FillSide {
-    LeftSide,
-    RightSide
-};
-
-} // namespace LoopBlinnConstants
-} // namespace WebCore
-
-#endif // LoopBlinnConstants_h
diff --git a/Source/WebCore/platform/graphics/gpu/LoopBlinnLocalTriangulator.cpp b/Source/WebCore/platform/graphics/gpu/LoopBlinnLocalTriangulator.cpp
deleted file mode 100644 (file)
index 1517a67..0000000
+++ /dev/null
@@ -1,279 +0,0 @@
-/*
- * Copyright (C) 2010 Google Inc. All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions
- * are met:
- *
- * 1.  Redistributions of source code must retain the above copyright
- *     notice, this list of conditions and the following disclaimer.
- * 2.  Redistributions in binary form must reproduce the above copyright
- *     notice, this list of conditions and the following disclaimer in the
- *     documentation and/or other materials provided with the distribution.
- *
- * THIS SOFTWARE IS PROVIDED BY APPLE AND ITS CONTRIBUTORS "AS IS" AND ANY
- * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
- * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
- * DISCLAIMED. IN NO EVENT SHALL APPLE OR ITS CONTRIBUTORS BE LIABLE FOR ANY
- * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
- * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
- * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
- * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF
- * THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- */
-
-#include "config.h"
-
-#if ENABLE(ACCELERATED_2D_CANVAS)
-
-#include "LoopBlinnLocalTriangulator.h"
-
-#include "LoopBlinnMathUtils.h"
-#include <algorithm>
-
-namespace WebCore {
-
-using LoopBlinnMathUtils::approxEqual;
-using LoopBlinnMathUtils::linesIntersect;
-using LoopBlinnMathUtils::pointInTriangle;
-
-bool LoopBlinnLocalTriangulator::Triangle::contains(LoopBlinnLocalTriangulator::Vertex* v)
-{
-    return indexForVertex(v) >= 0;
-}
-
-LoopBlinnLocalTriangulator::Vertex* LoopBlinnLocalTriangulator::Triangle::nextVertex(LoopBlinnLocalTriangulator::Vertex* current, bool traverseCounterClockwise)
-{
-    int index = indexForVertex(current);
-    ASSERT(index >= 0);
-    if (traverseCounterClockwise)
-        ++index;
-    else
-        --index;
-    if (index < 0)
-        index += 3;
-    else
-        index = index % 3;
-    return m_vertices[index];
-}
-
-int LoopBlinnLocalTriangulator::Triangle::indexForVertex(LoopBlinnLocalTriangulator::Vertex* vertex)
-{
-    for (int i = 0; i < 3; ++i)
-        if (m_vertices[i] == vertex)
-            return i;
-    return -1;
-}
-
-void LoopBlinnLocalTriangulator::Triangle::makeCounterClockwise()
-{
-    // Possibly swaps two vertices so that the triangle's vertices are
-    // always specified in counterclockwise order. This orders the
-    // vertices canonically when walking the interior edges from the
-    // start to the end vertex.
-    FloatPoint3D point0(m_vertices[0]->xyCoordinates());
-    FloatPoint3D point1(m_vertices[1]->xyCoordinates());
-    FloatPoint3D point2(m_vertices[2]->xyCoordinates());
-    FloatPoint3D crossProduct = (point1 - point0).cross(point2 - point0);
-    if (crossProduct.z() < 0)
-        std::swap(m_vertices[1], m_vertices[2]);
-}
-
-LoopBlinnLocalTriangulator::LoopBlinnLocalTriangulator()
-{
-    reset();
-}
-
-void LoopBlinnLocalTriangulator::reset()
-{
-    m_numberOfTriangles = 0;
-    m_numberOfInteriorVertices = 0;
-    for (int i = 0; i < 4; ++i) {
-        m_interiorVertices[i] = 0;
-        m_vertices[i].resetFlags();
-    }
-}
-
-void LoopBlinnLocalTriangulator::triangulate(InsideEdgeComputation computeInsideEdges, LoopBlinnConstants::FillSide sideToFill)
-{
-    triangulateHelper(sideToFill);
-
-    if (computeInsideEdges == ComputeInsideEdges) {
-        // We need to compute which vertices describe the path along the
-        // interior portion of the shape, to feed these vertices to the
-        // more general tessellation algorithm. It is possible that we
-        // could determine this directly while producing triangles above.
-        // Here we try to do it generally just by examining the triangles
-        // that have already been produced. We walk around them in a
-        // specific direction determined by which side of the curve is
-        // being filled. We ignore the interior vertex unless it is also
-        // the ending vertex, and skip the edges shared between two
-        // triangles.
-        Vertex* v = &m_vertices[0];
-        addInteriorVertex(v);
-        int numSteps = 0;
-        while (!v->end() && numSteps < 4) {
-            // Find the next vertex according to the above rules
-            bool gotNext = false;
-            for (int i = 0; i < numberOfTriangles() && !gotNext; ++i) {
-                Triangle* tri = getTriangle(i);
-                if (tri->contains(v)) {
-                    Vertex* next = tri->nextVertex(v, sideToFill == LoopBlinnConstants::RightSide);
-                    if (!next->marked() && !isSharedEdge(v, next) && (!next->interior() || next->end())) {
-                        addInteriorVertex(next);
-                        v = next;
-                        // Break out of for loop
-                        gotNext = true;
-                    }
-                }
-            }
-            ++numSteps;
-        }
-        if (!v->end()) {
-            // Something went wrong with the above algorithm; add the last
-            // vertex to the interior vertices anyway. (FIXME: should we
-            // add an assert here and do more extensive testing?)
-            addInteriorVertex(&m_vertices[3]);
-        }
-    }
-}
-
-void LoopBlinnLocalTriangulator::triangulateHelper(LoopBlinnConstants::FillSide sideToFill)
-{
-    reset();
-
-    m_vertices[3].setEnd(true);
-
-    // First test for degenerate cases.
-    for (int i = 0; i < 4; ++i) {
-        for (int j = i + 1; j < 4; ++j) {
-            if (approxEqual(m_vertices[i].xyCoordinates(), m_vertices[j].xyCoordinates())) {
-                // Two of the vertices are coincident, so we can eliminate at
-                // least one triangle. We might be able to eliminate the other
-                // as well, but this seems sufficient to avoid degenerate
-                // triangulations.
-                int indices[3] = { 0 };
-                int index = 0;
-                for (int k = 0; k < 4; ++k)
-                    if (k != j)
-                        indices[index++] = k;
-                addTriangle(&m_vertices[indices[0]],
-                            &m_vertices[indices[1]],
-                            &m_vertices[indices[2]]);
-                return;
-            }
-        }
-    }
-
-    // See whether any of the points are fully contained in the
-    // triangle defined by the other three.
-    for (int i = 0; i < 4; ++i) {
-        int indices[3] = { 0 };
-        int index = 0;
-        for (int j = 0; j < 4; ++j)
-            if (i != j)
-                indices[index++] = j;
-        if (pointInTriangle(m_vertices[i].xyCoordinates(),
-                            m_vertices[indices[0]].xyCoordinates(),
-                            m_vertices[indices[1]].xyCoordinates(),
-                            m_vertices[indices[2]].xyCoordinates())) {
-            // Produce three triangles surrounding this interior vertex.
-            for (int j = 0; j < 3; ++j)
-                addTriangle(&m_vertices[indices[j % 3]],
-                            &m_vertices[indices[(j + 1) % 3]],
-                            &m_vertices[i]);
-            // Mark the interior vertex so we ignore it if trying to trace
-            // the interior edge.
-            m_vertices[i].setInterior(true);
-            return;
-        }
-    }
-
-    // There are only a few permutations of the vertices, ignoring
-    // rotations, which are irrelevant:
-    //
-    //  0--3  0--2  0--3  0--1  0--2  0--1
-    //  |  |  |  |  |  |  |  |  |  |  |  |
-    //  |  |  |  |  |  |  |  |  |  |  |  |
-    //  1--2  1--3  2--1  2--3  3--1  3--2
-    //
-    // Note that three of these are reflections of each other.
-    // Therefore there are only three possible triangulations:
-    //
-    //  0--3  0--2  0--3
-    //  |\ |  |\ |  |\ |
-    //  | \|  | \|  | \|
-    //  1--2  1--3  2--1
-    //
-    // From which we can choose by seeing which of the potential
-    // diagonals intersect. Note that we choose the shortest diagonal
-    // to split the quad.
-    if (linesIntersect(m_vertices[0].xyCoordinates(),
-                       m_vertices[2].xyCoordinates(),
-                       m_vertices[1].xyCoordinates(),
-                       m_vertices[3].xyCoordinates())) {
-        if ((m_vertices[2].xyCoordinates() - m_vertices[0].xyCoordinates()).diagonalLengthSquared() <
-            (m_vertices[3].xyCoordinates() - m_vertices[1].xyCoordinates()).diagonalLengthSquared()) {
-            addTriangle(&m_vertices[0], &m_vertices[1], &m_vertices[2]);
-            addTriangle(&m_vertices[0], &m_vertices[2], &m_vertices[3]);
-        } else {
-            addTriangle(&m_vertices[0], &m_vertices[1], &m_vertices[3]);
-            addTriangle(&m_vertices[1], &m_vertices[2], &m_vertices[3]);
-        }
-    } else if (linesIntersect(m_vertices[0].xyCoordinates(),
-                              m_vertices[3].xyCoordinates(),
-                              m_vertices[1].xyCoordinates(),
-                              m_vertices[2].xyCoordinates())) {
-        if ((m_vertices[3].xyCoordinates() - m_vertices[0].xyCoordinates()).diagonalLengthSquared() <
-            (m_vertices[2].xyCoordinates() - m_vertices[1].xyCoordinates()).diagonalLengthSquared()) {
-            addTriangle(&m_vertices[0], &m_vertices[1], &m_vertices[3]);
-            addTriangle(&m_vertices[0], &m_vertices[3], &m_vertices[2]);
-        } else {
-            addTriangle(&m_vertices[0], &m_vertices[1], &m_vertices[2]);
-            addTriangle(&m_vertices[2], &m_vertices[1], &m_vertices[3]);
-        }
-    } else {
-        // Lines (0->1), (2->3) intersect -- or should, modulo numerical
-        // precision issues
-        if ((m_vertices[1].xyCoordinates() - m_vertices[0].xyCoordinates()).diagonalLengthSquared() <
-            (m_vertices[3].xyCoordinates() - m_vertices[2].xyCoordinates()).diagonalLengthSquared()) {
-            addTriangle(&m_vertices[0], &m_vertices[2], &m_vertices[1]);
-            addTriangle(&m_vertices[0], &m_vertices[1], &m_vertices[3]);
-        } else {
-            addTriangle(&m_vertices[0], &m_vertices[2], &m_vertices[3]);
-            addTriangle(&m_vertices[3], &m_vertices[2], &m_vertices[1]);
-        }
-    }
-}
-
-void LoopBlinnLocalTriangulator::addTriangle(Vertex* v0, Vertex* v1, Vertex* v2)
-{
-    ASSERT(m_numberOfTriangles < 3);
-    m_triangles[m_numberOfTriangles++].setVertices(v0, v1, v2);
-}
-
-void LoopBlinnLocalTriangulator::addInteriorVertex(Vertex* v)
-{
-    ASSERT(m_numberOfInteriorVertices < 4);
-    m_interiorVertices[m_numberOfInteriorVertices++] = v;
-    v->setMarked(true);
-}
-
-bool LoopBlinnLocalTriangulator::isSharedEdge(Vertex* v0, Vertex* v1)
-{
-    bool haveEdge01 = false;
-    bool haveEdge10 = false;
-    for (int i = 0; i < numberOfTriangles(); ++i) {
-        Triangle* tri = getTriangle(i);
-        if (tri->contains(v0) && tri->nextVertex(v0, true) == v1)
-            haveEdge01 = true;
-        if (tri->contains(v1) && tri->nextVertex(v1, true) == v0)
-            haveEdge10 = true;
-    }
-    return haveEdge01 && haveEdge10;
-}
-
-} // namespace WebCore
-
-#endif
diff --git a/Source/WebCore/platform/graphics/gpu/LoopBlinnLocalTriangulator.h b/Source/WebCore/platform/graphics/gpu/LoopBlinnLocalTriangulator.h
deleted file mode 100644 (file)
index aa2294e..0000000
+++ /dev/null
@@ -1,270 +0,0 @@
-/*
- * Copyright (C) 2010 Google Inc. All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions
- * are met:
- *
- * 1.  Redistributions of source code must retain the above copyright
- *     notice, this list of conditions and the following disclaimer.
- * 2.  Redistributions in binary form must reproduce the above copyright
- *     notice, this list of conditions and the following disclaimer in the
- *     documentation and/or other materials provided with the distribution.
- *
- * THIS SOFTWARE IS PROVIDED BY APPLE AND ITS CONTRIBUTORS "AS IS" AND ANY
- * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
- * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
- * DISCLAIMED. IN NO EVENT SHALL APPLE OR ITS CONTRIBUTORS BE LIABLE FOR ANY
- * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
- * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
- * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
- * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF
- * THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- */
-
-#ifndef LoopBlinnLocalTriangulator_h
-#define LoopBlinnLocalTriangulator_h
-
-#include "FloatPoint.h"
-#include "FloatPoint3D.h"
-#include "LoopBlinnConstants.h"
-#include <wtf/Assertions.h>
-#include <wtf/Noncopyable.h>
-
-namespace WebCore {
-
-// Performs a localized triangulation of the triangle mesh
-// corresponding to the four control point vertices of a cubic curve
-// segment.
-class LoopBlinnLocalTriangulator {
-    WTF_MAKE_NONCOPYABLE(LoopBlinnLocalTriangulator);
-public:
-    // The vertices that the triangulator operates upon, containing both
-    // the position information as well as the cubic texture
-    // coordinates.
-    class Vertex {
-        WTF_MAKE_NONCOPYABLE(Vertex);
-    public:
-        Vertex()
-        {
-            resetFlags();
-        }
-
-        const FloatPoint& xyCoordinates() const
-        {
-            return m_xyCoordinates;
-        }
-
-        const FloatPoint3D& klmCoordinates() const
-        {
-            return m_klmCoordinates;
-        }
-
-        // Sets the position and texture coordinates of the vertex.
-        void set(float x, float y,
-                 float k, float l, float m)
-        {
-            m_xyCoordinates.set(x, y);
-            m_klmCoordinates.set(k, l, m);
-        }
-
-        // Flags for walking from the start vertex to the end vertex.
-        bool end()
-        {
-            return m_end;
-        }
-
-        void setEnd(bool end)
-        {
-            m_end = end;
-        }
-
-        bool marked()
-        {
-            return m_marked;
-        }
-
-        void setMarked(bool marked)
-        {
-            m_marked = marked;
-        }
-
-        bool interior()
-        {
-            return m_interior;
-        }
-
-        void setInterior(bool interior)
-        {
-            m_interior = interior;
-        }
-
-        void resetFlags()
-        {
-            m_end = false;
-            m_marked = false;
-            m_interior = false;
-        }
-
-    private:
-        // 2D coordinates of the vertex in the plane.
-        FloatPoint m_xyCoordinates;
-        // Cubic texture coordinates for rendering the curve.
-        FloatPoint3D m_klmCoordinates;
-
-        // Flags for walking from the start vertex to the end vertex.
-        bool m_end;
-        bool m_marked;
-        bool m_interior;
-    };
-
-    // The triangles the Triangulator produces.
-    class Triangle {
-    public:
-        Triangle()
-        {
-            m_vertices[0] = 0;
-            m_vertices[1] = 0;
-            m_vertices[2] = 0;
-        }
-
-        // Gets the vertex at the given index, 0 <= index < 3.
-        Vertex* getVertex(int index)
-        {
-            ASSERT(index >= 0 && index < 3);
-            return m_vertices[index];
-        }
-
-        // Returns true if this triangle contains the given vertex (by
-        // identity, not geometrically).
-        bool contains(Vertex* v);
-
-        // Returns the vertex following the current one in the specified
-        // direction, counterclockwise or clockwise.
-        Vertex* nextVertex(Vertex* current, bool traverseCounterClockwise);
-
-        // Sets the vertices of this triangle, potentially reordering them
-        // to produce a canonical orientation.
-        void setVertices(Vertex* v0,
-                         Vertex* v1,
-                         Vertex* v2)
-        {
-            m_vertices[0] = v0;
-            m_vertices[1] = v1;
-            m_vertices[2] = v2;
-            makeCounterClockwise();
-        }
-
-    private:
-        // Returns the index [0..2] associated with the given vertex, or
-        // -1 if not found.
-        int indexForVertex(Vertex* vertex);
-
-        // Reorders the vertices in this triangle to make them
-        // counterclockwise when viewed in the 2D plane, in order to
-        // achieve a canonical ordering.
-        void makeCounterClockwise();
-
-        // Note: these are raw pointers because they point to the
-        // m_vertices contained in the surrounding triangulator.
-        Vertex* m_vertices[3];
-    };
-
-    LoopBlinnLocalTriangulator();
-
-    // Resets the triangulator's state. After each triangulation and
-    // before the next, call this to re-initialize the internal
-    // vertices' state.
-    void reset();
-
-    // Returns a mutable vertex stored in the triangulator. Use this to
-    // set up the vertices before a triangulation.
-    Vertex* getVertex(int index)
-    {
-        ASSERT(index >= 0 && index < 4);
-        return &m_vertices[index];
-    }
-
-    enum InsideEdgeComputation {
-        ComputeInsideEdges,
-        DontComputeInsideEdges
-    };
-
-    // Once the vertices' contents have been set up, call triangulate()
-    // to recompute the triangles.
-    //
-    // If computeInsideEdges is ComputeInsideEdges, then sideToFill
-    // will be used to determine which side of the cubic curve defined
-    // by the four control points is to be filled.
-    //
-    // The triangulation obeys the following guarantees:
-    //   - If the convex hull is a quadrilateral, then the shortest edge
-    //     will be chosen for the cut into two triangles.
-    //   - If one of the vertices is contained in the triangle spanned
-    //     by the other three, three triangles will be produced.
-    void triangulate(InsideEdgeComputation computeInsideEdges,
-                     LoopBlinnConstants::FillSide sideToFill);
-
-    // Number of triangles computed by triangulate().
-    int numberOfTriangles() const
-    {
-        return m_numberOfTriangles;
-    }
-
-    // Returns the computed triangle at index, 0 <= index < numberOfTriangles().
-    Triangle* getTriangle(int index)
-    {
-        ASSERT(index >= 0 && index < m_numberOfTriangles);
-        return &m_triangles[index];
-    }
-
-    // Number of vertices facing the inside of the shape, if
-    // ComputeInsideEdges was passed when triangulate() was called.
-    int numberOfInteriorVertices() const
-    {
-        return m_numberOfInteriorVertices;
-    }
-
-    // Fetches the given interior vertex, 0 <= index < numberOfInteriorVertices().
-    Vertex* getInteriorVertex(int index)
-    {
-        ASSERT(index >= 0 && index < m_numberOfInteriorVertices);
-        return m_interiorVertices[index];
-    }
-
-private:
-    void triangulateHelper(LoopBlinnConstants::FillSide sideToFill);
-
-    // Adds a triangle to the triangulation.
-    void addTriangle(Vertex* v0, Vertex* v1, Vertex* v2);
-
-    // Adds a vertex to the list of interior vertices.
-    void addInteriorVertex(Vertex* v);
-
-    // Indicates whether the edge between vertex v0 and v1 is shared
-    // between two or more triangles.
-    bool isSharedEdge(Vertex* v0, Vertex* v1);
-
-    // The vertices being triangulated.
-    Vertex m_vertices[4];
-
-    // The vertices corresponding to the edges facing the inside of the
-    // shape, in order from the start vertex to the end vertex. The more
-    // general triangulation algorithm tessellates this interior region.
-    Vertex* m_interiorVertices[4];
-    // The number of interior vertices that are valid for the current
-    // triangulation.
-    int m_numberOfInteriorVertices;
-
-    // There can be at most three triangles computed by this local
-    // algorithm, which occurs when one of the vertices is contained in
-    // the triangle spanned by the other three. Most of the time the
-    // algorithm computes two triangles.
-    Triangle m_triangles[3];
-    int m_numberOfTriangles;
-};
-
-} // namespace WebCore
-
-#endif // LoopBlinnLocalTriangulator_h
diff --git a/Source/WebCore/platform/graphics/gpu/LoopBlinnMathUtils.cpp b/Source/WebCore/platform/graphics/gpu/LoopBlinnMathUtils.cpp
deleted file mode 100644 (file)
index 87d72df..0000000
+++ /dev/null
@@ -1,666 +0,0 @@
-/*
- * Copyright (C) 2010 Google Inc. All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions
- * are met:
- *
- * 1.  Redistributions of source code must retain the above copyright
- *     notice, this list of conditions and the following disclaimer.
- * 2.  Redistributions in binary form must reproduce the above copyright
- *     notice, this list of conditions and the following disclaimer in the
- *     documentation and/or other materials provided with the distribution.
- *
- * THIS SOFTWARE IS PROVIDED BY APPLE AND ITS CONTRIBUTORS "AS IS" AND ANY
- * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
- * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
- * DISCLAIMED. IN NO EVENT SHALL APPLE OR ITS CONTRIBUTORS BE LIABLE FOR ANY
- * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
- * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
- * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
- * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF
- * THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- */
-
-#include "config.h"
-
-#include "LoopBlinnMathUtils.h"
-
-#include "FloatPoint.h"
-#include <algorithm>
-#include <wtf/MathExtras.h>
-
-namespace WebCore {
-namespace LoopBlinnMathUtils {
-
-namespace {
-
-// Utility functions local to this file.
-int orientation(const FloatPoint& p1,
-                const FloatPoint& p2,
-                const FloatPoint& p3)
-{
-    float crossProduct = (p2.y() - p1.y()) * (p3.x() - p2.x()) - (p3.y() - p2.y()) * (p2.x() - p1.x());
-    return (crossProduct < 0.0f) ? -1 : ((crossProduct > 0.0f) ? 1 : 0);
-}
-
-bool edgeEdgeTest(const FloatSize& v0Delta,
-                  const FloatPoint& v0,
-                  const FloatPoint& u0,
-                  const FloatPoint& u1)
-{
-    // This edge to edge test is based on Franlin Antonio's gem: "Faster
-    // Line Segment Intersection", in Graphics Gems III, pp. 199-202.
-    float ax = v0Delta.width();
-    float ay = v0Delta.height();
-    float bx = u0.x() - u1.x();
-    float by = u0.y() - u1.y();
-    float cx = v0.x() - u0.x();
-    float cy = v0.y() - u0.y();
-    float f = ay * bx - ax * by;
-    float d = by * cx - bx * cy;
-    if ((f > 0 && d >= 0 && d <= f) || (f < 0 && d <= 0 && d >= f)) {
-        float e = ax * cy - ay * cx;
-
-        // This additional test avoids reporting coincident edges, which
-        // is the behavior we want.
-        if (approxEqual(e, 0) || approxEqual(f, 0) || approxEqual(e, f))
-            return false;
-
-        if (f > 0)
-            return e >= 0 && e <= f;
-
-        return e <= 0 && e >= f;
-    }
-    return false;
-}
-
-bool edgeAgainstTriangleEdges(const FloatPoint& v0,
-                              const FloatPoint& v1,
-                              const FloatPoint& u0,
-                              const FloatPoint& u1,
-                              const FloatPoint& u2)
-{
-    FloatSize delta = v1 - v0;
-    // Test edge u0, u1 against v0, v1.
-    if (edgeEdgeTest(delta, v0, u0, u1))
-        return true;
-    // Test edge u1, u2 against v0, v1.
-    if (edgeEdgeTest(delta, v0, u1, u2))
-        return true;
-    // Test edge u2, u1 against v0, v1.
-    if (edgeEdgeTest(delta, v0, u2, u0))
-        return true;
-    return false;
-}
-
-// A roundoff factor in the cubic classification and texture coordinate
-// generation algorithms. It primarily determines the handling of corner
-// cases during the classification process. Be careful when adjusting it;
-// it has been determined empirically to work well. When changing it, you
-// should look in particular at shapes that contain quadratic curves and
-// ensure they still look smooth. Once pixel tests are running against this
-// algorithm, they should provide sufficient coverage to ensure that
-// adjusting the constant won't break anything.
-const float Epsilon = 5.0e-4f;
-
-} // anonymous namespace
-
-// Exported routines
-
-float roundToZero(float val)
-{
-    if (val < Epsilon && val > -Epsilon)
-        return 0;
-    return val;
-}
-
-bool approxEqual(const FloatPoint& v0, const FloatPoint& v1)
-{
-    return (v0 - v1).diagonalLengthSquared() < Epsilon * Epsilon;
-}
-
-bool approxEqual(const FloatPoint3D& v0, const FloatPoint3D& v1)
-{
-    return (v0 - v1).lengthSquared() < Epsilon * Epsilon;
-}
-
-bool approxEqual(float f0, float f1)
-{
-    return fabsf(f0 - f1) < Epsilon;
-}
-
-bool linesIntersect(const FloatPoint& p1,
-                    const FloatPoint& q1,
-                    const FloatPoint& p2,
-                    const FloatPoint& q2)
-{
-    return (orientation(p1, q1, p2) != orientation(p1, q1, q2)
-            && orientation(p2, q2, p1) != orientation(p2, q2, q1));
-}
-
-bool pointInTriangle(const FloatPoint& point,
-                     const FloatPoint& a,
-                     const FloatPoint& b,
-                     const FloatPoint& c)
-{
-    // Algorithm from http://www.blackpawn.com/texts/pointinpoly/default.html
-    float x0 = c.x() - a.x();
-    float y0 = c.y() - a.y();
-    float x1 = b.x() - a.x();
-    float y1 = b.y() - a.y();
-    float x2 = point.x() - a.x();
-    float y2 = point.y() - a.y();
-
-    float dot00 = x0 * x0 + y0 * y0;
-    float dot01 = x0 * x1 + y0 * y1;
-    float dot02 = x0 * x2 + y0 * y2;
-    float dot11 = x1 * x1 + y1 * y1;
-    float dot12 = x1 * x2 + y1 * y2;
-    float denominator = dot00 * dot11 - dot01 * dot01;
-    if (!denominator)
-        // Triangle is zero-area. Treat query point as not being inside.
-        return false;
-    // Compute
-    float inverseDenominator = 1.0f / denominator;
-    float u = (dot11 * dot02 - dot01 * dot12) * inverseDenominator;
-    float v = (dot00 * dot12 - dot01 * dot02) * inverseDenominator;
-
-    return (u > 0.0f) && (v > 0.0f) && (u + v < 1.0f);
-}
-
-bool trianglesOverlap(const FloatPoint& a1,
-                      const FloatPoint& b1,
-                      const FloatPoint& c1,
-                      const FloatPoint& a2,
-                      const FloatPoint& b2,
-                      const FloatPoint& c2)
-{
-    // Derived from coplanar_tri_tri() at
-    // http://jgt.akpeters.com/papers/ShenHengTang03/tri_tri.html ,
-    // simplified for the 2D case and modified so that overlapping edges
-    // do not report overlapping triangles.
-
-    // Test all edges of triangle 1 against the edges of triangle 2.
-    if (edgeAgainstTriangleEdges(a1, b1, a2, b2, c2)
-        || edgeAgainstTriangleEdges(b1, c1, a2, b2, c2)
-        || edgeAgainstTriangleEdges(c1, a1, a2, b2, c2))
-        return true;
-    // Finally, test if tri1 is totally contained in tri2 or vice versa.
-    // The paper above only performs the first two point-in-triangle tests.
-    // Because we define that triangles sharing a vertex or edge don't
-    // overlap, we must perform additional tests to see whether one
-    // triangle is contained in the other.
-    if (pointInTriangle(a1, a2, b2, c2)
-        || pointInTriangle(a2, a1, b1, c1)
-        || pointInTriangle(b1, a2, b2, c2)
-        || pointInTriangle(b2, a1, b1, c1)
-        || pointInTriangle(c1, a2, b2, c2)
-        || pointInTriangle(c2, a1, b1, c1))
-        return true;
-    return false;
-}
-
-namespace {
-
-// Helper routines for public XRay queries below. All of this code
-// originated in Skia; see include/core/ and src/core/, SkScalar.h and
-// SkGeometry.{cpp,h}.
-
-const float NearlyZeroConstant = (1.0f / (1 << 12));
-
-bool nearlyZero(float x, float tolerance = NearlyZeroConstant)
-{
-    ASSERT(tolerance > 0.0f);
-    return ::fabsf(x) < tolerance;
-}
-
-// Linearly interpolate between a and b, based on t.
-// If t is 0, return a; if t is 1, return b; else interpolate.
-// t must be [0..1].
-float interpolate(float a, float b, float t)
-{
-    ASSERT(t >= 0 && t <= 1);
-    return a + (b - a) * t;
-}
-
-float evaluateCubic(float controlPoint0, float controlPoint1, float controlPoint2, float controlPoint3, float t)
-{
-    ASSERT(t >= 0 && t <= 1);
-
-    if (!t)
-        return controlPoint0;
-
-    float ab = interpolate(controlPoint0, controlPoint1, t);
-    float bc = interpolate(controlPoint1, controlPoint2, t);
-    float cd = interpolate(controlPoint2, controlPoint3, t);
-    float abc = interpolate(ab, bc, t);
-    float bcd = interpolate(bc, cd, t);
-    return interpolate(abc, bcd, t);
-}
-
-// Evaluates the point on the source cubic specified by t, 0 <= t <= 1.0.
-FloatPoint evaluateCubicAt(const FloatPoint cubic[4], float t)
-{
-    return FloatPoint(evaluateCubic(cubic[0].x(), cubic[1].x(), cubic[2].x(), cubic[3].x(), t),
-                      evaluateCubic(cubic[0].y(), cubic[1].y(), cubic[2].y(), cubic[3].y(), t));
-}
-
-bool xRayCrossesMonotonicCubic(const XRay& xRay, const FloatPoint cubic[4], bool& ambiguous)
-{
-    ambiguous = false;
-
-    // Find the minimum and maximum y of the extrema, which are the
-    // first and last points since this cubic is monotonic
-    float minY = std::min(cubic[0].y(), cubic[3].y());
-    float maxY = std::max(cubic[0].y(), cubic[3].y());
-
-    if (xRay.y() == cubic[0].y()
-        || xRay.y() < minY
-        || xRay.y() > maxY) {
-        // The query line definitely does not cross the curve
-        ambiguous = (xRay.y() == cubic[0].y());
-        return false;
-    }
-
-    const bool pointAtExtremum = (xRay.y() == cubic[3].y());
-
-    float minX = std::min(std::min(std::min(cubic[0].x(), cubic[1].x()),
-                                   cubic[2].x()),
-                          cubic[3].x());
-    if (xRay.x() < minX) {
-        // The query line definitely crosses the curve
-        ambiguous = pointAtExtremum;
-        return true;
-    }
-
-    float maxX = std::max(std::max(std::max(cubic[0].x(), cubic[1].x()),
-                                   cubic[2].x()),
-                          cubic[3].x());
-    if (xRay.x() > maxX)
-        // The query line definitely does not cross the curve
-        return false;
-
-    // Do a binary search to find the parameter value which makes y as
-    // close as possible to the query point. See whether the query
-    // line's origin is to the left of the associated x coordinate.
-
-    // MaxIterations is chosen as the number of mantissa bits for a float,
-    // since there's no way we are going to get more precision by
-    // iterating more times than that.
-    const int MaxIterations = 23;
-    FloatPoint evaluatedPoint;
-    int iter = 0;
-    float upperT;
-    float lowerT;
-    // Need to invert direction of t parameter if cubic goes up
-    // instead of down
-    if (cubic[3].y() > cubic[0].y()) {
-        upperT = 1;
-        lowerT = 0;
-    } else {
-        upperT = 0;
-        lowerT = 1;
-    }
-    do {
-        float t = 0.5f * (upperT + lowerT);
-        evaluatedPoint = evaluateCubicAt(cubic, t);
-        if (xRay.y() > evaluatedPoint.y())
-            lowerT = t;
-        else
-            upperT = t;
-    } while (++iter < MaxIterations && !nearlyZero(evaluatedPoint.y() - xRay.y()));
-
-    // FIXME: once we have more regression tests for this code,
-    // determine whether this should be using a fuzzy test.
-    if (xRay.x() <= evaluatedPoint.x()) {
-        ambiguous = pointAtExtremum;
-        return true;
-    }
-    return false;
-}
-
-// Divides the numerator by the denominator safely for the case where
-// the result must lie in the range (0..1). Result indicates whether
-// the result is valid.
-bool safeUnitDivide(float numerator, float denominator, float& ratio)
-{
-    if (numerator < 0) {
-        // Make the "numerator >= denominator" check below work.
-        numerator = -numerator;
-        denominator = -denominator;
-    }
-    if (!numerator || !denominator || numerator >= denominator)
-        return false;
-    float r = numerator / denominator;
-    if (std::isnan(r))
-        return false;
-    ASSERT(r >= 0 && r < 1);
-    if (!r) // catch underflow if numerator <<<< denominator
-        return false;
-    ratio = r;
-    return true;
-}
-
-// From Numerical Recipes in C.
-//
-//   q = -1/2 (b + sign(b) sqrt[b*b - 4*a*c])
-//   x1 = q / a
-//   x2 = c / q
-//
-// Returns the number of real roots of the equation [0..2]. Roots are
-// returned in sorted order, smaller root first.
-int findUnitQuadRoots(float a, float b, float c, float roots[2])
-{
-    if (!a)
-        return safeUnitDivide(-c, b, roots[0]) ? 1 : 0;
-
-    float discriminant = b*b - 4*a*c;
-    if (discriminant < 0 || std::isnan(discriminant)) // complex roots
-        return 0;
-    discriminant = sqrtf(discriminant);
-
-    float q = (b < 0) ? -(b - discriminant) / 2 : -(b + discriminant) / 2;
-    int numberOfRoots = 0;
-    if (safeUnitDivide(q, a, roots[numberOfRoots]))
-        ++numberOfRoots;
-    if (safeUnitDivide(c, q, roots[numberOfRoots]))
-        ++numberOfRoots;
-    if (numberOfRoots == 2) {
-        // Seemingly have two roots. Check for equality and sort.
-        if (roots[0] == roots[1])
-            return 1;
-        if (roots[0] > roots[1])
-            std::swap(roots[0], roots[1]);
-    }
-    return numberOfRoots;
-}
-
-// Cubic'(t) = pt^2 + qt + r, where
-//   p = 3(-a + 3(b - c) + d)
-//   q = 6(a - 2b + c)
-//   r = 3(b - a)
-// Solve for t, keeping only those that fit between 0 < t < 1.
-int findCubicExtrema(float a, float b, float c, float d, float tValues[2])
-{
-    // Divide p, q, and r by 3 to simplify the equations.
-    float p = d - a + 3*(b - c);
-    float q = 2*(a - b - b + c);
-    float r = b - a;
-
-    return findUnitQuadRoots(p, q, r, tValues);
-}
-
-void interpolateCubicCoords(float controlPoint0, float controlPoint1, float controlPoint2, float controlPoint3, float* dst, float t)
-{
-    float ab = interpolate(controlPoint0, controlPoint1, t);
-    float bc = interpolate(controlPoint1, controlPoint2, t);
-    float cd = interpolate(controlPoint2, controlPoint3, t);
-    float abc = interpolate(ab, bc, t);
-    float bcd = interpolate(bc, cd, t);
-    float abcd = interpolate(abc, bcd, t);
-
-    dst[0] = controlPoint0;
-    dst[2] = ab;
-    dst[4] = abc;
-    dst[6] = abcd;
-    dst[8] = bcd;
-    dst[10] = cd;
-    dst[12] = controlPoint3;
-}
-
-#ifndef NDEBUG
-bool isUnitInterval(float x)
-{
-    return x > 0 && x < 1;
-}
-#endif
-
-void chopCubicAtTValues(const FloatPoint src[4], FloatPoint dst[], const float tValues[], int roots)
-{
-#ifndef NDEBUG
-    for (int i = 0; i < roots - 1; ++i) {
-        ASSERT(isUnitInterval(tValues[i]));
-        ASSERT(isUnitInterval(tValues[i+1]));
-        ASSERT(tValues[i] < tValues[i+1]);
-    }
-#endif
-
-    if (!roots) {
-        // nothing to chop
-        for (int j = 0; j < 4; ++j)
-            dst[j] = src[j];
-        return;
-    }
-
-    float t = tValues[0];
-    FloatPoint tmp[4];
-    for (int j = 0; j < 4; ++j)
-        tmp[j] = src[j];
-
-    for (int i = 0; i < roots; ++i) {
-        chopCubicAt(tmp, dst, t);
-        if (i == roots - 1)
-            break;
-
-        dst += 3;
-        // Make tmp contain the remaining cubic (after the first chop).
-        for (int j = 0; j < 4; ++j)
-            tmp[j] = dst[j];
-
-        // Watch out for the case that the renormalized t isn't in range.
-        if (!safeUnitDivide(tValues[i+1] - tValues[i], 1.0f - tValues[i], t)) {
-            // If it isn't, just create a degenerate cubic.
-            dst[4] = dst[5] = dst[6] = tmp[3];
-            break;
-        }
-    }
-}
-
-void flattenDoubleCubicYExtrema(FloatPoint coords[7])
-{
-    coords[2].setY(coords[3].y());
-    coords[4].setY(coords[3].y());
-}
-
-int chopCubicAtYExtrema(const FloatPoint src[4], FloatPoint dst[10])
-{
-    float tValues[2];
-    int roots = findCubicExtrema(src[0].y(), src[1].y(), src[2].y(), src[3].y(), tValues);
-
-    chopCubicAtTValues(src, dst, tValues, roots);
-    if (roots) {
-        // we do some cleanup to ensure our Y extrema are flat
-        flattenDoubleCubicYExtrema(&dst[0]);
-        if (roots == 2)
-            flattenDoubleCubicYExtrema(&dst[3]);
-    }
-    return roots;
-}
-
-} // anonymous namespace
-
-// Public cubic operations.
-
-void chopCubicAt(const FloatPoint src[4], FloatPoint dst[7], float t)
-{
-    ASSERT(t >= 0 && t <= 1);
-
-    float output[14];
-    interpolateCubicCoords(src[0].x(), src[1].x(), src[2].x(), src[3].x(), &output[0], t);
-    interpolateCubicCoords(src[0].y(), src[1].y(), src[2].y(), src[3].y(), &output[1], t);
-    for (int i = 0; i < 7; i++)
-        dst[i].set(output[2 * i], output[2 * i + 1]);
-}
-
-// Public XRay queries.
-
-bool xRayCrossesLine(const XRay& xRay, const FloatPoint pts[2], bool& ambiguous)
-{
-    ambiguous = false;
-
-    // Determine quick discards.
-    // Consider query line going exactly through point 0 to not
-    // intersect, for symmetry with xRayCrossesMonotonicCubic.
-    if (xRay.y() == pts[0].y()) {
-        ambiguous = true;
-        return false;
-    }
-    if (xRay.y() < pts[0].y() && xRay.y() < pts[1].y())
-        return false;
-    if (xRay.y() > pts[0].y() && xRay.y() > pts[1].y())
-        return false;
-    if (xRay.x() > pts[0].x() && xRay.x() > pts[1].x())
-        return false;
-    // Determine degenerate cases
-    if (nearlyZero(pts[0].y() - pts[1].y()))
-        return false;
-    if (nearlyZero(pts[0].x() - pts[1].x())) {
-        // We've already determined the query point lies within the
-        // vertical range of the line segment.
-        if (xRay.x() <= pts[0].x()) {
-            ambiguous = (xRay.y() == pts[1].y());
-            return true;
-        }
-        return false;
-    }
-    // Ambiguity check
-    if (xRay.y() == pts[1].y()) {
-        if (xRay.x() <= pts[1].x()) {
-            ambiguous = true;
-            return true;
-        }
-        return false;
-    }
-    // Full line segment evaluation
-    float deltaY = pts[1].y() - pts[0].y();
-    float deltaX = pts[1].x() - pts[0].x();
-    float slope = deltaY / deltaX;
-    float b = pts[0].y() - slope * pts[0].x();
-    // Solve for x coordinate at y = xRay.y()
-    float x = (xRay.y() - b) / slope;
-    return xRay.x() <= x;
-}
-
-int numXRayCrossingsForCubic(const XRay& xRay, const FloatPoint cubic[4], bool& ambiguous)
-{
-    int numCrossings = 0;
-    FloatPoint monotonicCubics[10];
-    int numMonotonicCubics = 1 + chopCubicAtYExtrema(cubic, monotonicCubics);
-    ambiguous = false;
-    FloatPoint* monotonicCubicsPointer = &monotonicCubics[0];
-    for (int i = 0; i < numMonotonicCubics; ++i) {
-        if (xRayCrossesMonotonicCubic(xRay, monotonicCubicsPointer, ambiguous))
-            ++numCrossings;
-        if (ambiguous)
-            return 0;
-        monotonicCubicsPointer += 3;
-    }
-    return numCrossings;
-}
-
-/*
- * Based on C code from the article
- * "Testing the Convexity of a Polygon"
- * by Peter Schorn and Frederick Fisher,
- * (schorn@inf.ethz.ch, fred@kpc.com)
- * in "Graphics Gems IV", Academic Press, 1994
- */
-
-static inline int convexCompare(const FloatSize& delta)
-{
-    return (delta.width() > 0) ? -1 : /* x coord diff, second pt > first pt */
-           (delta.width() < 0) ?  1 : /* x coord diff, second pt < first pt */
-           (delta.height() > 0) ? -1 : /* x coord same, second pt > first pt */
-           (delta.height() < 0) ?  1 : /* x coord same, second pt > first pt */
-           0; /* second pt equals first point */
-}
-
-static inline float convexCross(const FloatSize& p, const FloatSize& q)
-{
-    return p.width() * q.height() - p.height() * q.width();
-}
-
-static inline bool convexCheckTriple(const FloatSize& dcur, const FloatSize& dprev, int* curDir, int* dirChanges, int* angleSign)
-{
-    int thisDir = convexCompare(dcur);
-    if (thisDir == -*curDir)
-        ++*dirChanges;
-    *curDir = thisDir;
-    float cross = convexCross(dprev, dcur);
-    if (cross > 0) {
-        if (*angleSign == -1)
-            return false;
-        *angleSign = 1;
-    } else if (cross < 0) {
-        if (*angleSign == 1)
-            return false;
-        *angleSign = -1;
-    }
-    return true;
-}
-
-bool isConvex(const FloatPoint* vertices, int nVertices)
-{
-    int dirChanges = 0, angleSign = 0;
-    FloatPoint second, third;
-    FloatSize dprev, dcur;
-
-    /* Get different point, return if less than 3 diff points. */
-    if (nVertices < 3)
-        return false;
-    int i = 1;
-    while (true) {
-        second = vertices[i++];
-        dprev = second - vertices[0];
-        if (dprev.width() || dprev.height())
-            break;
-        /* Check if out of points. Check here to avoid slowing down cases
-         * without repeated points.
-         */
-        if (i >= nVertices)
-            return false;
-    }
-    FloatPoint saveSecond = second;
-    int curDir = convexCompare(dprev);        /* Find initial direction */
-    while (i < nVertices) {
-        /* Get different point, break if no more points */
-        third = vertices[i++];
-        dcur = third - second;
-        if (!dcur.width() && !dcur.height())
-            continue;
-
-        /* Check current three points */
-        if (!convexCheckTriple(dcur, dprev, &curDir, &dirChanges, &angleSign)) 
-            return false;
-        second = third;     /* Remember ptr to current point. */
-        dprev = dcur;       /* Remember current delta. */
-    }
-
-    /* Must check for direction changes from last vertex back to first */
-    third = vertices[0];                  /* Prepare for 'ConvexCheckTriple' */
-    dcur = third - second;
-    if (convexCompare(dcur)) {
-        if (!convexCheckTriple(dcur, dprev, &curDir, &dirChanges, &angleSign)) 
-            return false;
-        second = third;     /* Remember ptr to current point. */
-        dprev = dcur;       /* Remember current delta. */
-    }
-
-    /* and check for direction changes back to second vertex */
-    dcur = saveSecond - second;
-    if (!convexCheckTriple(dcur, dprev, &curDir, &dirChanges, &angleSign)) 
-        return false;
-
-    /* Decide on polygon type given accumulated status */
-    if (dirChanges > 2)
-        return false;
-
-    if (angleSign > 0 || angleSign < 0)
-        return true;
-    return false;
-}
-
-} // namespace LoopBlinnMathUtils
-} // namespace WebCore
diff --git a/Source/WebCore/platform/graphics/gpu/LoopBlinnMathUtils.h b/Source/WebCore/platform/graphics/gpu/LoopBlinnMathUtils.h
deleted file mode 100644 (file)
index 361d901..0000000
+++ /dev/null
@@ -1,111 +0,0 @@
-/*
- * Copyright (C) 2010 Google Inc. All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions
- * are met:
- *
- * 1.  Redistributions of source code must retain the above copyright
- *     notice, this list of conditions and the following disclaimer.
- * 2.  Redistributions in binary form must reproduce the above copyright
- *     notice, this list of conditions and the following disclaimer in the
- *     documentation and/or other materials provided with the distribution.
- *
- * THIS SOFTWARE IS PROVIDED BY APPLE AND ITS CONTRIBUTORS "AS IS" AND ANY
- * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
- * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
- * DISCLAIMED. IN NO EVENT SHALL APPLE OR ITS CONTRIBUTORS BE LIABLE FOR ANY
- * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
- * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
- * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
- * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF
- * THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- */
-
-#ifndef LoopBlinnMathUtils_h
-#define LoopBlinnMathUtils_h
-
-#include "FloatPoint.h"
-#include "FloatPoint3D.h"
-#include <math.h>
-
-namespace WebCore {
-
-// Use a namespace for these so we can easily import them.
-namespace LoopBlinnMathUtils {
-
-float roundToZero(float val);
-bool approxEqual(const FloatPoint& v0, const FloatPoint& v1);
-bool approxEqual(const FloatPoint3D& v0, const FloatPoint3D& v1);
-bool approxEqual(float f0, float f1);
-
-// Determines whether the line segment between (p1, q1) intersects
-// that between (p2, q2).
-bool linesIntersect(const FloatPoint& p1,
-                    const FloatPoint& q1,
-                    const FloatPoint& p2,
-                    const FloatPoint& q2);
-
-// Determines whether "point" is inside the 2D triangle defined by
-// vertices a, b, and c. This test defines that points exactly on an
-// edge are not considered to be inside the triangle.
-bool pointInTriangle(const FloatPoint& point,
-                     const FloatPoint& a,
-                     const FloatPoint& b,
-                     const FloatPoint& c);
-
-// Determines whether the triangles defined by the points (a1, b1, c1)
-// and (a2, b2, c2) overlap. The definition of this function is that
-// if the two triangles only share an adjacent edge or vertex, they
-// are not considered to overlap.
-bool trianglesOverlap(const FloatPoint& a1,
-                      const FloatPoint& b1,
-                      const FloatPoint& c1,
-                      const FloatPoint& a2,
-                      const FloatPoint& b2,
-                      const FloatPoint& c2);
-
-// Given a src cubic bezier, chops it at the specified t value,
-// where 0 < t < 1, and returns the two new cubics in dst[0..3]
-// and dst[3..6].
-void chopCubicAt(const FloatPoint src[4], FloatPoint dst[7], float t);
-
-// "X-Ray" queries. An XRay is a half-line originating at the given
-// point and extending to x=+infinity.
-typedef FloatPoint XRay;
-
-// Given an arbitrary cubic bezier, return the number of times an XRay
-// crosses the cubic. Valid return values are [0..3].
-//
-// By definition the cubic is open at the starting point; in other
-// words, if pt.fY is equivalent to cubic[0].fY, and pt.fX is to the
-// left of the curve, the line is not considered to cross the curve,
-// but if it is equal to cubic[3].fY then it is considered to
-// cross.
-//
-// Outgoing "ambiguous" argument indicates whether the answer is ambiguous
-// because the query occurred exactly at one of the endpoints' y
-// coordinates or at a tangent point, indicating that another query y
-// coordinate is preferred for robustness.
-int numXRayCrossingsForCubic(const XRay& xRay,
-                             const FloatPoint cubic[4],
-                             bool& ambiguous);
-
-// Given a line segment from lineEndpoints[0] to lineEndpoints[1], and an
-// XRay, returns true if they intersect. Outgoing "ambiguous" argument
-// indicates whether the answer is ambiguous because the query occurred
-// exactly at one of the endpoints' y coordinates, indicating that another
-// query y coordinate is preferred for robustness.
-bool xRayCrossesLine(const XRay& xRay,
-                     const FloatPoint lineEndpoints[2],
-                     bool& ambiguous);
-
-
-bool isConvex(const FloatPoint* vertices, int nVertices);
-
-} // namespace LoopBlinnMathUtils
-
-} // namespace WebCore
-
-#endif // LoopBlinnMathUtils_h
diff --git a/Source/WebCore/platform/graphics/gpu/LoopBlinnPathCache.cpp b/Source/WebCore/platform/graphics/gpu/LoopBlinnPathCache.cpp
deleted file mode 100644 (file)
index 35f15e5..0000000
+++ /dev/null
@@ -1,86 +0,0 @@
-/*
- * Copyright (C) 2011 Google Inc. All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions
- * are met:
- *
- * 1.  Redistributions of source code must retain the above copyright
- *     notice, this list of conditions and the following disclaimer.
- * 2.  Redistributions in binary form must reproduce the above copyright
- *     notice, this list of conditions and the following disclaimer in the
- *     documentation and/or other materials provided with the distribution.
- *
- * THIS SOFTWARE IS PROVIDED BY APPLE AND ITS CONTRIBUTORS "AS IS" AND ANY
- * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
- * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
- * DISCLAIMED. IN NO EVENT SHALL APPLE OR ITS CONTRIBUTORS BE LIABLE FOR ANY
- * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
- * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
- * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
- * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF
- * THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- */
-
-#include "config.h"
-
-#include "LoopBlinnPathCache.h"
-
-namespace WebCore {
-
-LoopBlinnPathCache::LoopBlinnPathCache()
-{
-}
-
-LoopBlinnPathCache::~LoopBlinnPathCache()
-{
-}
-
-void LoopBlinnPathCache::addVertex(float x, float y,
-                                   float k, float l, float m)
-{
-    m_vertices.append(x);
-    m_vertices.append(y);
-    m_texcoords.append(k);
-    m_texcoords.append(l);
-    m_texcoords.append(m);
-}
-
-void LoopBlinnPathCache::clear()
-{
-    m_vertices.clear();
-    m_texcoords.clear();
-    m_interiorVertices.clear();
-#ifdef LOOP_BLINN_PATH_CACHE_DEBUG_INTERIOR_EDGES
-    m_interiorEdgeVertices.clear();
-#endif // LOOP_BLINN_PATH_CACHE_DEBUG_INTERIOR_EDGES
-}
-
-void LoopBlinnPathCache::addInteriorVertex(float x, float y)
-{
-    m_interiorVertices.append(x);
-    m_interiorVertices.append(y);
-}
-
-#ifdef LOOP_BLINN_PATH_CACHE_DEBUG_INTERIOR_EDGES
-unsigned LoopBlinnPathCache::numberOfInteriorEdgeVertices() const
-{
-    return m_interiorEdgeVertices.size() / 2;
-}
-
-const float* LoopBlinnPathCache::interiorEdgeVertices() const
-{
-    if (!numberOfInteriorEdgeVertices())
-        return 0;
-    return m_interiorEdgeVertices.data();
-}
-
-void LoopBlinnPathCache::addInteriorEdgeVertex(float x, float y)
-{
-    m_interiorEdgeVertices.append(x);
-    m_interiorEdgeVertices.append(y);
-}
-#endif // LOOP_BLINN_PATH_CACHE_DEBUG_INTERIOR_EDGES
-
-} // namespace WebCore
diff --git a/Source/WebCore/platform/graphics/gpu/LoopBlinnPathCache.h b/Source/WebCore/platform/graphics/gpu/LoopBlinnPathCache.h
deleted file mode 100644 (file)
index d21d246..0000000
+++ /dev/null
@@ -1,123 +0,0 @@
-/*
- * Copyright (C) 2011 Google Inc. All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions
- * are met:
- *
- * 1.  Redistributions of source code must retain the above copyright
- *     notice, this list of conditions and the following disclaimer.
- * 2.  Redistributions in binary form must reproduce the above copyright
- *     notice, this list of conditions and the following disclaimer in the
- *     documentation and/or other materials provided with the distribution.
- *
- * THIS SOFTWARE IS PROVIDED BY APPLE AND ITS CONTRIBUTORS "AS IS" AND ANY
- * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
- * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
- * DISCLAIMED. IN NO EVENT SHALL APPLE OR ITS CONTRIBUTORS BE LIABLE FOR ANY
- * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
- * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
- * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
- * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF
- * THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- */
-
-#ifndef LoopBlinnPathCache_h
-#define LoopBlinnPathCache_h
-
-#include <wtf/Noncopyable.h>
-#include <wtf/Vector.h>
-
-namespace WebCore {
-
-// A cache of the processed triangle mesh for a given path. Because these
-// might be expensive to allocate (using malloc/free internally), it is
-// recommended to try to reuse them when possible.
-
-// Uncomment the following to obtain debugging information for the edges
-// facing the interior region of the mesh.
-// #define LOOP_BLINN_PATH_CACHE_DEBUG_INTERIOR_EDGES
-
-class LoopBlinnPathCache {
-    WTF_MAKE_NONCOPYABLE(LoopBlinnPathCache);
-public:
-    LoopBlinnPathCache();
-    ~LoopBlinnPathCache();
-
-    unsigned numberOfVertices() const { return m_vertices.size() / 2; }
-
-    // Get the base pointer to the vertex information. There are two
-    // coordinates per vertex. This pointer is valid until the cache is
-    // cleared or another vertex is added. Returns 0 if there are no
-    // vertices in the mesh.
-    const float* vertices() const
-    {
-        if (!numberOfVertices())
-            return 0;
-        return m_vertices.data();
-    }
-
-    // Get the base pointer to the texture coordinate information. There
-    // are three coordinates per vertex. This pointer is valid until the
-    // cache is cleared or another vertex is added. Returns 0 if
-    // there are no vertices in the mesh.
-    const float* texcoords() const
-    {
-        if (!numberOfVertices())
-            return 0;
-        return m_texcoords.data();
-    }
-
-    // Adds a vertex's information to the cache. The first two arguments
-    // are the x and y coordinates of the vertex on the plane; the last
-    // three arguments are the cubic texture coordinates associated with
-    // this vertex.
-    void addVertex(float x, float y,
-                   float /*k*/, float /*l*/, float /*m*/);
-
-    unsigned numberOfInteriorVertices() const { return m_interiorVertices.size() / 2; }
-
-    // Base pointer to the interior vertices; two coordinates per
-    // vertex, which can be drawn as GL_TRIANGLES. Returns 0 if there
-    // are no interior vertices in the mesh.
-    const float* interiorVertices() const
-    {
-        if (!numberOfInteriorVertices())
-            return 0;
-        return m_interiorVertices.data();
-    }
-
-    void addInteriorVertex(float x, float y);
-
-    // Clears all of the stored vertex information in this cache.
-    void clear();
-
-#ifdef LOOP_BLINN_PATH_CACHE_DEBUG_INTERIOR_EDGES
-    // The number of interior edge vertices
-    unsigned numberOfInteriorEdgeVertices() const;
-    // Base pointer to the interior vertices; two coordinates per
-    // vertex, which can be drawn as GL_LINES. Returns 0 if there are
-    // no interior edge vertices in the mesh.
-    const float* interiorEdgeVertices() const;
-    void addInteriorEdgeVertex(float x, float y);
-#endif // LOOP_BLINN_PATH_CACHE_DEBUG_INTERIOR_EDGES
-
-private:
-    // The two-dimensional vertices of the triangle mesh.
-    Vector<float> m_vertices;
-
-    // The three-dimensional cubic texture coordinates.
-    Vector<float> m_texcoords;
-
-    Vector<float> m_interiorVertices;
-
-#ifdef LOOP_BLINN_PATH_CACHE_DEBUG_INTERIOR_EDGES
-    // The following is only for debugging
-    Vector<float> m_interiorEdgeVertices;
-#endif // LOOP_BLINN_PATH_CACHE_DEBUG_INTERIOR_EDGES
-};
-
-} // namespace WebCore
-
-#endif // LoopBlinnPathCache_h
diff --git a/Source/WebCore/platform/graphics/gpu/LoopBlinnPathProcessor.h b/Source/WebCore/platform/graphics/gpu/LoopBlinnPathProcessor.h
deleted file mode 100644 (file)
index ad89bc1..0000000
+++ /dev/null
@@ -1,126 +0,0 @@
-/*
- * Copyright (C) 2011 Google Inc. All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions
- * are met:
- *
- * 1.  Redistributions of source code must retain the above copyright
- *     notice, this list of conditions and the following disclaimer.
- * 2.  Redistributions in binary form must reproduce the above copyright
- *     notice, this list of conditions and the following disclaimer in the
- *     documentation and/or other materials provided with the distribution.
- *
- * THIS SOFTWARE IS PROVIDED BY APPLE AND ITS CONTRIBUTORS "AS IS" AND ANY
- * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
- * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
- * DISCLAIMED. IN NO EVENT SHALL APPLE OR ITS CONTRIBUTORS BE LIABLE FOR ANY
- * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
- * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
- * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
- * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF
- * THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- */
-
-// The main entry point for Loop and Blinn's GPU accelerated curve
-// rendering algorithm.
-
-#ifndef LoopBlinnPathProcessor_h
-#define LoopBlinnPathProcessor_h
-
-#include <wtf/Noncopyable.h>
-#include <wtf/PassRefPtr.h>
-#include <wtf/RefPtr.h>
-#include <wtf/Vector.h>
-
-namespace WebCore {
-
-// We use a namespace for classes which are simply implementation
-// details of the algorithm but which we need to reference from the
-// class definition.
-namespace LoopBlinnPathProcessorImplementation {
-
-class Contour;
-class Segment;
-
-} // namespace LoopBlinnPathProcessorImplementation
-
-class Path;
-class LoopBlinnPathCache;
-class PODArena;
-
-// The LoopBlinnPathProcessor turns a Path (assumed to contain one or
-// more closed regions) into a set of exterior and interior triangles,
-// stored in the LoopBlinnPathCache. The exterior triangles have
-// associated 3D texture coordinates which are used to evaluate the
-// curve's inside/outside function on a per-pixel basis. The interior
-// triangles are filled with 100% opacity.
-//
-// Note that the fill style and management of multiple layers are
-// separate concerns, handled at a higher level with shaders and
-// polygon offsets.
-class LoopBlinnPathProcessor {
-public:
-    LoopBlinnPathProcessor();
-    explicit LoopBlinnPathProcessor(PassRefPtr<PODArena>);
-    ~LoopBlinnPathProcessor();
-
-    // Transforms the given path into a triangle mesh for rendering
-    // using Loop and Blinn's shader, placing the result into the given
-    // LoopBlinnPathCache.
-    void process(const Path&, LoopBlinnPathCache&);
-
-#ifndef NDEBUG
-    // Enables or disables verbose logging in debug mode.
-    void setVerboseLogging(bool onOrOff);
-#endif
-
-private:
-    // Builds a list of contours for the given path.
-    void buildContours(const Path&);
-
-    // Determines whether the left or right side of each contour should
-    // be filled.
-    void determineSidesToFill();
-
-    // Determines whether the given (closed) contour is oriented
-    // clockwise or counterclockwise.
-    void determineOrientation(LoopBlinnPathProcessorImplementation::Contour*);
-
-    // Subdivides the curves so that there are no overlaps of the
-    // triangles associated with the curves' control points.
-    void subdivideCurves();
-
-    // Helper function used during curve subdivision.
-    void conditionallySubdivide(LoopBlinnPathProcessorImplementation::Segment*,
-                                Vector<LoopBlinnPathProcessorImplementation::Segment*>& nextSegments);
-
-    // Tessellates the interior regions of the contours.
-    void tessellateInterior(LoopBlinnPathCache&);
-
-#ifndef NDEBUG
-    // For debugging the orientation computation. Returns all of the
-    // segments overlapping the given Y coordinate.
-    Vector<LoopBlinnPathProcessorImplementation::Segment*> allSegmentsOverlappingY(LoopBlinnPathProcessorImplementation::Contour*, float x, float y);
-
-    // For debugging the curve subdivision algorithm. Subdivides the
-    // curves using an alternate, slow (O(n^3)) algorithm.
-    void subdivideCurvesSlow();
-#endif
-
-    // PODArena from which to allocate temporary objects.
-    RefPtr<PODArena> m_arena;
-
-    // The contours described by the path.
-    Vector<LoopBlinnPathProcessorImplementation::Contour*> m_contours;
-
-#ifndef NDEBUG
-    // Whether or not to perform verbose logging in debug mode.
-    bool m_verboseLogging;
-#endif
-};
-
-} // namespace WebCore
-
-#endif // LoopBlinnPathProcessor_h
diff --git a/Source/WebCore/platform/graphics/gpu/LoopBlinnShader.cpp b/Source/WebCore/platform/graphics/gpu/LoopBlinnShader.cpp
deleted file mode 100644 (file)
index 5d9fdc8..0000000
+++ /dev/null
@@ -1,59 +0,0 @@
-/*
- * Copyright (C) 2011 Google Inc. All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions
- * are met:
- *
- * 1.  Redistributions of source code must retain the above copyright
- *     notice, this list of conditions and the following disclaimer.
- * 2.  Redistributions in binary form must reproduce the above copyright
- *     notice, this list of conditions and the following disclaimer in the
- *     documentation and/or other materials provided with the distribution.
- *
- * THIS SOFTWARE IS PROVIDED BY APPLE AND ITS CONTRIBUTORS "AS IS" AND ANY
- * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
- * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
- * DISCLAIMED. IN NO EVENT SHALL APPLE OR ITS CONTRIBUTORS BE LIABLE FOR ANY
- * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
- * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
- * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
- * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF
- * THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- */
-
-#include "config.h"
-
-#include "LoopBlinnShader.h"
-
-#include "GraphicsContext3D.h"
-
-namespace WebCore {
-
-LoopBlinnShader::LoopBlinnShader(GraphicsContext3D* context, unsigned program)
-    : Shader(context, program)
-{
-    m_worldViewProjectionLocation = context->getUniformLocation(program, "worldViewProjection");
-    m_positionLocation = context->getAttribLocation(program, "position");
-    m_klmLocation = context->getAttribLocation(program, "klm");
-}
-
-void LoopBlinnShader::use(unsigned vertexOffset, unsigned klmOffset, const AffineTransform& transform)
-{
-    m_context->useProgram(m_program);
-
-    float matrix[16];
-    affineTo4x4(transform, matrix);
-    m_context->uniformMatrix4fv(m_worldViewProjectionLocation, 1 /*count*/, false /*transpose*/, matrix);
-
-    m_context->vertexAttribPointer(m_positionLocation, 2, GraphicsContext3D::FLOAT, false, 0, vertexOffset);
-    m_context->enableVertexAttribArray(m_positionLocation);
-
-    if (m_klmLocation != -1) {
-        m_context->vertexAttribPointer(m_klmLocation, 3, GraphicsContext3D::FLOAT, false, 0, klmOffset);
-        m_context->enableVertexAttribArray(m_klmLocation);
-    }
-}
-
-} // namespace WebCore
diff --git a/Source/WebCore/platform/graphics/gpu/LoopBlinnShader.h b/Source/WebCore/platform/graphics/gpu/LoopBlinnShader.h
deleted file mode 100644 (file)
index 2d24dc3..0000000
+++ /dev/null
@@ -1,58 +0,0 @@
-/*
- * Copyright (C) 2011 Google Inc. All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions
- * are met:
- *
- * 1.  Redistributions of source code must retain the above copyright
- *     notice, this list of conditions and the following disclaimer.
- * 2.  Redistributions in binary form must reproduce the above copyright
- *     notice, this list of conditions and the following disclaimer in the
- *     documentation and/or other materials provided with the distribution.
- *
- * THIS SOFTWARE IS PROVIDED BY APPLE AND ITS CONTRIBUTORS "AS IS" AND ANY
- * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
- * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
- * DISCLAIMED. IN NO EVENT SHALL APPLE OR ITS CONTRIBUTORS BE LIABLE FOR ANY
- * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
- * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
- * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
- * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF
- * THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- */
-
-#ifndef LoopBlinnShader_h
-#define LoopBlinnShader_h
-
-#include "Shader.h"
-
-namespace WebCore {
-
-class GraphicsContext3D;
-
-class LoopBlinnShader : public Shader {
-public:
-    enum Region {
-        Interior,
-        Exterior
-    };
-
-protected:
-    LoopBlinnShader(GraphicsContext3D*, unsigned program);
-
-    // This assumes the vertices and klm coordinates are stored in the
-    // same, currently bound, buffer object, contiguously and at the
-    // specified offsets.
-    void use(unsigned vertexOffset, unsigned klmOffset, const AffineTransform&);
-
-private:
-    int m_worldViewProjectionLocation;
-    int m_positionLocation;
-    int m_klmLocation;
-};
-
-} // namespace WebCore
-
-#endif // LoopBlinnShader_h
diff --git a/Source/WebCore/platform/graphics/gpu/LoopBlinnSolidFillShader.cpp b/Source/WebCore/platform/graphics/gpu/LoopBlinnSolidFillShader.cpp
deleted file mode 100644 (file)
index 80c8677..0000000
+++ /dev/null
@@ -1,63 +0,0 @@
-/*
- * Copyright (C) 2011 Google Inc. All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions
- * are met:
- *
- * 1.  Redistributions of source code must retain the above copyright
- *     notice, this list of conditions and the following disclaimer.
- * 2.  Redistributions in binary form must reproduce the above copyright
- *     notice, this list of conditions and the following disclaimer in the
- *     documentation and/or other materials provided with the distribution.
- *
- * THIS SOFTWARE IS PROVIDED BY APPLE AND ITS CONTRIBUTORS "AS IS" AND ANY
- * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
- * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
- * DISCLAIMED. IN NO EVENT SHALL APPLE OR ITS CONTRIBUTORS BE LIABLE FOR ANY
- * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
- * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
- * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
- * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF
- * THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- */
-
-#include "config.h"
-
-#include "LoopBlinnSolidFillShader.h"
-
-#include "Color.h"
-#include "GraphicsContext3D.h"
-
-namespace WebCore {
-
-PassOwnPtr<LoopBlinnSolidFillShader> LoopBlinnSolidFillShader::create(GraphicsContext3D* context,
-                                                                      LoopBlinnShader::Region region,
-                                                                      Shader::AntialiasType antialiasType)
-{
-    VertexType type = (region == Interior) ? LoopBlinnInterior : LoopBlinnExterior;
-    unsigned program = loadProgram(context,
-                                   generateVertex(type, SolidFill),
-                                   generateFragment(type, SolidFill, antialiasType));
-    if (!program)
-        return nullptr;
-    return adoptPtr(new LoopBlinnSolidFillShader(context, program));
-}
-
-LoopBlinnSolidFillShader::LoopBlinnSolidFillShader(GraphicsContext3D* context, unsigned program)
-    : LoopBlinnShader(context, program)
-{
-    m_colorLocation = context->getUniformLocation(program, "color");
-}
-
-void LoopBlinnSolidFillShader::use(unsigned vertexOffset, unsigned klmOffset, const AffineTransform& transform, const Color& color)
-{
-    LoopBlinnShader::use(vertexOffset, klmOffset, transform);
-
-    float rgba[4];
-    color.getRGBA(rgba[0], rgba[1], rgba[2], rgba[3]);
-    m_context->uniform4f(m_colorLocation, rgba[0] * rgba[3], rgba[1] * rgba[3], rgba[2] * rgba[3], rgba[3]);
-}
-
-} // namespace WebCore
diff --git a/Source/WebCore/platform/graphics/gpu/LoopBlinnSolidFillShader.h b/Source/WebCore/platform/graphics/gpu/LoopBlinnSolidFillShader.h
deleted file mode 100644 (file)
index 36312a2..0000000
+++ /dev/null
@@ -1,52 +0,0 @@
-/*
- * Copyright (C) 2011 Google Inc. All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions
- * are met:
- *
- * 1.  Redistributions of source code must retain the above copyright
- *     notice, this list of conditions and the following disclaimer.
- * 2.  Redistributions in binary form must reproduce the above copyright
- *     notice, this list of conditions and the following disclaimer in the
- *     documentation and/or other materials provided with the distribution.
- *
- * THIS SOFTWARE IS PROVIDED BY APPLE AND ITS CONTRIBUTORS "AS IS" AND ANY
- * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
- * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
- * DISCLAIMED. IN NO EVENT SHALL APPLE OR ITS CONTRIBUTORS BE LIABLE FOR ANY
- * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
- * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
- * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
- * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF
- * THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- */
-
-#ifndef LoopBlinnSolidFillShader_h
-#define LoopBlinnSolidFillShader_h
-
-#include "LoopBlinnShader.h"
-
-namespace WebCore {
-
-class GraphicsContext3D;
-
-class LoopBlinnSolidFillShader : public LoopBlinnShader {
-public:
-    static PassOwnPtr<LoopBlinnSolidFillShader> create(GraphicsContext3D*, Region, AntialiasType);
-
-    // This assumes the vertices and klm coordinates are stored in the
-    // same, currently bound, buffer object, contiguously and at the
-    // specified offsets.
-    void use(unsigned vertexOffset, unsigned klmOffset, const AffineTransform&, const Color&);
-
-private:
-    LoopBlinnSolidFillShader(GraphicsContext3D*, unsigned program);
-
-    int m_colorLocation;
-};
-
-} // namespace WebCore
-
-#endif // LoopBlinnSolidFillShader_h
diff --git a/Source/WebCore/platform/graphics/gpu/LoopBlinnTextureCoords.cpp b/Source/WebCore/platform/graphics/gpu/LoopBlinnTextureCoords.cpp
deleted file mode 100644 (file)
index d272fe1..0000000
+++ /dev/null
@@ -1,175 +0,0 @@
-/*
- * Copyright (C) 2010 Google Inc. All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions
- * are met:
- *
- * 1.  Redistributions of source code must retain the above copyright
- *     notice, this list of conditions and the following disclaimer.
- * 2.  Redistributions in binary form must reproduce the above copyright
- *     notice, this list of conditions and the following disclaimer in the
- *     documentation and/or other materials provided with the distribution.
- *
- * THIS SOFTWARE IS PROVIDED BY APPLE AND ITS CONTRIBUTORS "AS IS" AND ANY
- * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
- * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
- * DISCLAIMED. IN NO EVENT SHALL APPLE OR ITS CONTRIBUTORS BE LIABLE FOR ANY
- * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
- * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
- * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
- * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF
- * THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- */
-
-#include "config.h"
-
-#if ENABLE(ACCELERATED_2D_CANVAS)
-
-#include "LoopBlinnTextureCoords.h"
-
-#include <math.h>
-#include <wtf/Assertions.h>
-
-namespace WebCore {
-
-LoopBlinnTextureCoords::Result LoopBlinnTextureCoords::compute(const LoopBlinnClassifier::Result& classification, LoopBlinnConstants::FillSide sideToFill)
-{
-    // Loop and Blinn's formulation states that the right side of the
-    // curve is defined to be the inside (filled region), but for some
-    // reason it looks like with the default orientation parameters we
-    // are filling the left side of the curve. Regardless, because we
-    // can receive arbitrarily oriented curves as input, we might have
-    // to reverse the orientation of the cubic texture coordinates even
-    // in cases where the paper doesn't say it is necessary.
-    bool reverseOrientation = false;
-    static const float OneThird = 1.0f / 3.0f;
-    static const float TwoThirds = 2.0f / 3.0f;
-    LoopBlinnClassifier::CurveType curveType = classification.curveType;
-
-    LoopBlinnTextureCoords::Result result;
-
-    switch (curveType) {
-    case LoopBlinnClassifier::kSerpentine: {
-        float t1 = sqrtf(9.0f * classification.d2 * classification.d2 - 12 * classification.d1 * classification.d3);
-        float ls = 3.0f * classification.d2 - t1;
-        float lt = 6.0f * classification.d1;
-        float ms = 3.0f * classification.d2 + t1;
-        float mt = lt;
-        float ltMinusLs = lt - ls;
-        float mtMinusMs = mt - ms;
-        result.klmCoordinates[0] = FloatPoint3D(ls * ms,
-                                                ls * ls * ls,
-                                                ms * ms * ms);
-        result.klmCoordinates[1] = FloatPoint3D(OneThird * (3.0f * ls * ms - ls * mt - lt * ms),
-                                                ls * ls * (ls - lt),
-                                                ms * ms * (ms - mt));
-        result.klmCoordinates[2] = FloatPoint3D(OneThird * (lt * (mt - 2.0f * ms) + ls * (3.0f * ms - 2.0f * mt)),
-                                                ltMinusLs * ltMinusLs * ls,
-                                                mtMinusMs * mtMinusMs * ms);
-        result.klmCoordinates[3] = FloatPoint3D(ltMinusLs * mtMinusMs,
-                                                -(ltMinusLs * ltMinusLs * ltMinusLs),
-                                                -(mtMinusMs * mtMinusMs * mtMinusMs));
-        if (classification.d1 < 0.0f)
-            reverseOrientation = true;
-        break;
-    }
-
-    case LoopBlinnClassifier::kLoop: {
-        float t1 = sqrtf(4.0f * classification.d1 * classification.d3 - 3.0f * classification.d2 * classification.d2);
-        float ls = classification.d2 - t1;
-        float lt = 2.0f * classification.d1;
-        float ms = classification.d2 + t1;
-        float mt = lt;
-
-        // Figure out whether there is a rendering artifact requiring
-        // the curve to be subdivided by the caller.
-        float ql = ls / lt;
-        float qm = ms / mt;
-        if (0.0f < ql && ql < 1.0f) {
-            result.hasRenderingArtifact = true;
-            result.subdivisionParameterValue = ql;
-            return result;
-        }
-
-        if (0.0f < qm && qm < 1.0f) {
-            result.hasRenderingArtifact = true;
-            result.subdivisionParameterValue = qm;
-            return result;
-        }
-
-        float ltMinusLs = lt - ls;
-        float mtMinusMs = mt - ms;
-        result.klmCoordinates[0] = FloatPoint3D(ls * ms,
-                                                ls * ls * ms,
-                                                ls * ms * ms);
-        result.klmCoordinates[1] = FloatPoint3D(OneThird * (-ls * mt - lt * ms + 3.0f * ls * ms),
-                                                -OneThird * ls * (ls * (mt - 3.0f * ms) + 2.0f * lt * ms),
-                                                -OneThird * ms * (ls * (2.0f * mt - 3.0f * ms) + lt * ms));
-        result.klmCoordinates[2] = FloatPoint3D(OneThird * (lt * (mt - 2.0f * ms) + ls * (3.0f * ms - 2.0f * mt)),
-                                                OneThird * (lt - ls) * (ls * (2.0f * mt - 3.0f * ms) + lt * ms),
-                                                OneThird * (mt - ms) * (ls * (mt - 3.0f * ms) + 2.0f * lt * ms));
-        result.klmCoordinates[3] = FloatPoint3D(ltMinusLs * mtMinusMs,
-                                                -(ltMinusLs * ltMinusLs) * mtMinusMs,
-                                                -ltMinusLs * mtMinusMs * mtMinusMs);
-        reverseOrientation = ((classification.d1 > 0.0f && result.klmCoordinates[0].x() < 0.0f)
-                           || (classification.d1 < 0.0f && result.klmCoordinates[0].x() > 0.0f));
-        break;
-    }
-
-    case LoopBlinnClassifier::kCusp: {
-        float ls = classification.d3;
-        float lt = 3.0f * classification.d2;
-        float lsMinusLt = ls - lt;
-        result.klmCoordinates[0] = FloatPoint3D(ls,
-                                                ls * ls * ls,
-                                                1.0f);
-        result.klmCoordinates[1] = FloatPoint3D(ls - OneThird * lt,
-                                                ls * ls * lsMinusLt,
-                                                1.0f);
-        result.klmCoordinates[2] = FloatPoint3D(ls - TwoThirds * lt,
-                                                lsMinusLt * lsMinusLt * ls,
-                                                1.0f);
-        result.klmCoordinates[3] = FloatPoint3D(lsMinusLt,
-                                                lsMinusLt * lsMinusLt * lsMinusLt,
-                                                1.0f);
-        break;
-    }
-
-    case LoopBlinnClassifier::kQuadratic: {
-        result.klmCoordinates[0] = FloatPoint3D(0, 0, 0);
-        result.klmCoordinates[1] = FloatPoint3D(OneThird, 0, OneThird);
-        result.klmCoordinates[2] = FloatPoint3D(TwoThirds, OneThird, TwoThirds);
-        result.klmCoordinates[3] = FloatPoint3D(1, 1, 1);
-        if (classification.d3 < 0)
-            reverseOrientation = true;
-        break;
-    }
-
-    case LoopBlinnClassifier::kLine:
-    case LoopBlinnClassifier::kPoint:
-        result.isLineOrPoint = true;
-        break;
-
-    default:
-        ASSERT_NOT_REACHED();
-        break;
-    }
-
-    if (sideToFill == LoopBlinnConstants::RightSide)
-        reverseOrientation = !reverseOrientation;
-
-    if (reverseOrientation) {
-        for (int i = 0; i < 4; ++i) {
-            result.klmCoordinates[i].setX(-result.klmCoordinates[i].x());
-            result.klmCoordinates[i].setY(-result.klmCoordinates[i].y());
-        }
-    }
-
-    return result;
-}
-
-} // namespace WebCore
-
-#endif
diff --git a/Source/WebCore/platform/graphics/gpu/LoopBlinnTextureCoords.h b/Source/WebCore/platform/graphics/gpu/LoopBlinnTextureCoords.h
deleted file mode 100644 (file)
index 5fdeb3b..0000000
+++ /dev/null
@@ -1,82 +0,0 @@
-/*
- * Copyright (C) 2010 Google Inc. All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions
- * are met:
- *
- * 1.  Redistributions of source code must retain the above copyright
- *     notice, this list of conditions and the following disclaimer.
- * 2.  Redistributions in binary form must reproduce the above copyright
- *     notice, this list of conditions and the following disclaimer in the
- *     documentation and/or other materials provided with the distribution.
- *
- * THIS SOFTWARE IS PROVIDED BY APPLE AND ITS CONTRIBUTORS "AS IS" AND ANY
- * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
- * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
- * DISCLAIMED. IN NO EVENT SHALL APPLE OR ITS CONTRIBUTORS BE LIABLE FOR ANY
- * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
- * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
- * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
- * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF
- * THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- */
-
-#ifndef LoopBlinnTextureCoords_h
-#define LoopBlinnTextureCoords_h
-
-#include "FloatPoint3D.h"
-#include "LoopBlinnClassifier.h"
-#include "LoopBlinnConstants.h"
-
-#include <wtf/Noncopyable.h>
-
-namespace WebCore {
-
-// Computes three-dimensional texture coordinates for the control
-// points of a cubic curve for rendering via the shader in "Rendering
-// Vector Art on the GPU" by Loop and Blinn, GPU Gems 3, Chapter 25.
-class LoopBlinnTextureCoords {
-public:
-    // Container for the cubic texture coordinates and other associated
-    // information.
-    struct Result {
-        Result()
-            : isLineOrPoint(false)
-            , hasRenderingArtifact(false)
-            , subdivisionParameterValue(0.0f) { }
-
-        // The (k, l, m) texture coordinates that are to be associated
-        // with the four control points of the cubic curve.
-        FloatPoint3D klmCoordinates[4];
-
-        // Indicates whether the curve is a line or a point, in which case
-        // we do not need to add its triangles to the mesh.
-        bool isLineOrPoint;
-
-        // For the loop case, indicates whether a rendering artifact was
-        // detected, in which case the curve needs to be further
-        // subdivided.
-        bool hasRenderingArtifact;
-
-        // If a rendering artifact will occur for the given loop curve,
-        // this is the parameter value (0 <= value <= 1) at which the
-        // curve needs to be subdivided to fix the artifact.
-        float subdivisionParameterValue;
-    };
-
-    // Computes the texture coordinates for a cubic curve segment's
-    // control points, given the classification of the curve as well as
-    // an indication of which side is to be filled.
-    static Result compute(const LoopBlinnClassifier::Result& classification,
-                          LoopBlinnConstants::FillSide sideToFill);
-
-private:
-    // This class does not need to be instantiated.
-    LoopBlinnTextureCoords() { }
-};
-
-} // namespace WebCore
-
-#endif // LoopBlinnTextureCoords_h
diff --git a/Source/WebCore/platform/graphics/gpu/Shader.cpp b/Source/WebCore/platform/graphics/gpu/Shader.cpp
deleted file mode 100644 (file)
index 13c5ebf..0000000
+++ /dev/null
@@ -1,277 +0,0 @@
-/*
- * Copyright (c) 2010, Google Inc. All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions are
- * met:
- *
- *     * Redistributions of source code must retain the above copyright
- * notice, this list of conditions and the following disclaimer.
- *     * Redistributions in binary form must reproduce the above
- * copyright notice, this list of conditions and the following disclaimer
- * in the documentation and/or other materials provided with the
- * distribution.
- *     * Neither the name of Google Inc. nor the names of its
- * contributors may be used to endorse or promote products derived from
- * this software without specific prior written permission.
- *
- * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
- * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
- * LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR
- * A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT
- * OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
- * SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT
- * LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
- * DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY
- * THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
- * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- */
-
-#include "config.h"
-
-#if ENABLE(ACCELERATED_2D_CANVAS)
-
-#include "Shader.h"
-
-#include "AffineTransform.h"
-#include "GraphicsContext3D.h"
-
-#include <wtf/text/CString.h>
-#include <wtf/text/StringBuilder.h>
-
-namespace WebCore {
-
-// static
-void Shader::affineTo3x3(const AffineTransform& transform, float mat[9])
-{
-    mat[0] = transform.a();
-    mat[1] = transform.b();
-    mat[2] = 0.0f;
-    mat[3] = transform.c();
-    mat[4] = transform.d();
-    mat[5] = 0.0f;
-    mat[6] = transform.e();
-    mat[7] = transform.f();
-    mat[8] = 1.0f;
-}
-
-// static
-void Shader::affineTo4x4(const AffineTransform& transform, float mat[16])
-{
-    mat[0] = transform.a();
-    mat[1] = transform.b();
-    mat[2] = 0.0f;
-    mat[3] = 0.0f;
-    mat[4] = transform.c();
-    mat[5] = transform.d();
-    mat[6] = 0.0f;
-    mat[7] = 0.0f;
-    mat[8] = 0.0f;
-    mat[9] = 0.0f;
-    mat[10] = 1.0f;
-    mat[11] = 0.0f;
-    mat[12] = transform.e();
-    mat[13] = transform.f();
-    mat[14] = 0.0f;
-    mat[15] = 1.0f;
-}
-
-// static
-unsigned Shader::loadShader(GraphicsContext3D* context, unsigned type, const String& shaderSource)
-{
-    unsigned shader = context->createShader(type);
-    if (!shader)
-        return 0;
-
-    context->shaderSource(shader, shaderSource);
-    context->compileShader(shader);
-    int compileStatus = 0;
-    context->getShaderiv(shader, GraphicsContext3D::COMPILE_STATUS, &compileStatus);
-    if (!compileStatus) {
-        String infoLog = context->getShaderInfoLog(shader);
-        LOG_ERROR("%s", infoLog.utf8().data());
-        context->deleteShader(shader);
-        return 0;
-    }
-    return shader;
-}
-
-// static
-unsigned Shader::loadProgram(GraphicsContext3D* context, const String& vertexShaderSource, const String& fragmentShaderSource)
-{
-    unsigned vertexShader = loadShader(context, GraphicsContext3D::VERTEX_SHADER, vertexShaderSource);
-    if (!vertexShader)
-        return 0;
-    unsigned fragmentShader = loadShader(context, GraphicsContext3D::FRAGMENT_SHADER, fragmentShaderSource);
-    if (!fragmentShader)
-        return 0;
-    unsigned program = context->createProgram();
-    if (!program)
-        return 0;
-    context->attachShader(program, vertexShader);
-    context->attachShader(program, fragmentShader);
-    context->linkProgram(program);
-    int linkStatus = 0;
-    context->getProgramiv(program, GraphicsContext3D::LINK_STATUS, &linkStatus);
-    if (!linkStatus)
-        context->deleteProgram(program);
-    context->deleteShader(vertexShader);
-    context->deleteShader(fragmentShader);
-    return program;
-}
-
-Shader::Shader(GraphicsContext3D* context, unsigned program)
-    : m_context(context)
-    , m_program(program)
-{
-}
-
-Shader::~Shader()
-{
-    m_context->deleteProgram(m_program);
-}
-
-// static
-String Shader::generateVertex(Shader::VertexType vertexType, Shader::FillType fillType)
-{
-    StringBuilder builder;
-    switch (vertexType) {
-    case TwoDimensional:
-        builder.append(
-            "uniform mat3 matrix;\n"
-            "attribute vec2 position;\n");
-        break;
-    case LoopBlinnInterior:
-        builder.append(
-            "uniform mat4 worldViewProjection;\n"
-            "attribute vec2 position;\n");
-        break;
-    case LoopBlinnExterior:
-        builder.append(
-            "uniform mat4 worldViewProjection;\n"
-            "attribute vec2 position;\n"
-            "attribute vec3 klm;\n"
-            "varying vec3 v_klm;\n");
-        break;
-    }
-
-    if (fillType == TextureFill) {
-        builder.append(
-            "uniform mat3 texMatrix;\n"
-            "varying vec3 texCoord;\n");
-    }
-
-    builder.append(
-        "void main() {\n");
-
-    if (vertexType == TwoDimensional) {
-        builder.append(
-            "gl_Position = vec4(matrix * vec3(position, 1.0), 1.0);\n");
-    } else {
-        builder.append(
-            "gl_Position = worldViewProjection * vec4(position, 0.0, 1.0);\n");
-        if (vertexType == LoopBlinnExterior) {
-            builder.append(
-                "v_klm = klm;\n");
-        }
-    }
-
-    if (fillType == TextureFill) {
-        builder.append(
-            "texCoord = texMatrix * vec3(position, 1.0);\n");
-    }
-
-    builder.append(
-        "}\n");
-
-    return builder.toString();
-}
-
-// static
-String Shader::generateFragment(Shader::VertexType vertexType, Shader::FillType fillType, Shader::AntialiasType antialiasType)
-{
-    StringBuilder builder;
-    builder.append(
-        "#ifdef GL_ES\n"
-        "precision mediump float;\n"
-        "#endif\n");
-
-    if (vertexType == LoopBlinnExterior) {
-        if (antialiasType == Antialiased) {
-            builder.append(
-                "#extension GL_OES_standard_derivatives : enable\n");
-        }
-        builder.append(
-            "varying vec3 v_klm;\n");
-    }
-
-    switch (fillType) {
-    case SolidFill:
-        builder.append(
-            "uniform vec4 color;\n");
-        break;
-    case TextureFill:
-        builder.append(
-            "uniform sampler2D sampler;\n"
-            "uniform float globalAlpha;\n"
-            "varying vec3 texCoord;\n");
-        break;
-    }
-
-    builder.append(
-        "void main() {\n");
-
-    if (vertexType != LoopBlinnExterior) {
-        builder.append(
-            "float alpha = 1.0;\n");
-    } else {
-        if (antialiasType == Antialiased) {
-            builder.append(
-                "  // Gradients\n"
-                "  vec3 px = dFdx(v_klm);\n"
-                "  vec3 py = dFdy(v_klm);\n"
-                "\n"
-                "  // Chain rule\n"
-                "  float k2 = v_klm.x * v_klm.x;\n"
-                "  float c = k2 * v_klm.x - v_klm.y * v_klm.z;\n"
-                "  float k23 = 3.0 * k2;\n"
-                "  float cx = k23 * px.x - v_klm.z * px.y - v_klm.y * px.z;\n"
-                "  float cy = k23 * py.x - v_klm.z * py.y - v_klm.y * py.z;\n"
-                "\n"
-                "  // Signed distance\n"
-                "  float sd = c / sqrt(cx * cx + cy * cy);\n"
-                "\n"
-                "  // Linear alpha\n"
-                "  // FIXME: figure out why this needs to be\n"
-                "  // negated compared to the HLSL version, and also why\n"
-                "  // we need an adjustment by +1.0 for it to look good.\n"
-                "  // float alpha = clamp(0.5 - sd, 0.0, 1.0);\n"
-                "  float alpha = clamp(sd + 0.5, 0.0, 1.0);\n");
-        } else {
-            builder.append(
-                "  float t = v_klm.x * v_klm.x * v_klm.x - v_klm.y * v_klm.z;\n"
-                "  float alpha = clamp(sign(t), 0.0, 1.0);\n");
-        }
-    }
-
-    switch (fillType) {
-    case SolidFill:
-        builder.append(
-            "gl_FragColor = color * alpha;\n");
-        break;
-    case TextureFill:
-        builder.append(
-            "gl_FragColor = texture2D(sampler, texCoord.xy) * alpha * globalAlpha;\n");
-        break;
-    }
-
-    builder.append(
-        "}\n");
-
-    return builder.toString();
-}
-
-} // namespace WebCore
-
-#endif
diff --git a/Source/WebCore/platform/graphics/gpu/Shader.h b/Source/WebCore/platform/graphics/gpu/Shader.h
deleted file mode 100644 (file)
index 35d1a3b..0000000
+++ /dev/null
@@ -1,82 +0,0 @@
-/*
- * Copyright (c) 2010, Google Inc. All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions are
- * met:
- *
- *     * Redistributions of source code must retain the above copyright
- * notice, this list of conditions and the following disclaimer.
- *     * Redistributions in binary form must reproduce the above
- * copyright notice, this list of conditions and the following disclaimer
- * in the documentation and/or other materials provided with the
- * distribution.
- *     * Neither the name of Google Inc. nor the names of its
- * contributors may be used to endorse or promote products derived from
- * this software without specific prior written permission.
- *
- * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
- * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
- * LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR
- * A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT
- * OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
- * SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT
- * LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
- * DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY
- * THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
- * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- */
-
-#ifndef Shader_h
-#define Shader_h
-
-#include <wtf/Noncopyable.h>
-#include <wtf/PassOwnPtr.h>
-#include <wtf/text/WTFString.h>
-
-namespace WebCore {
-
-class AffineTransform;
-class GraphicsContext3D;
-class Color;
-
-class Shader {
-    WTF_MAKE_NONCOPYABLE(Shader);
-public:
-    enum VertexType {
-        TwoDimensional,
-        LoopBlinnInterior,
-        LoopBlinnExterior
-    };
-
-    enum FillType {
-        SolidFill,
-        TextureFill
-    };
-
-    // Currently only applies to the Loop-Blinn vertex type.
-    enum AntialiasType {
-        NotAntialiased,
-        Antialiased
-    };
-
-protected:
-    Shader(GraphicsContext3D*, unsigned program);
-    ~Shader();
-
-    static String generateVertex(VertexType, FillType);
-    static String generateFragment(VertexType, FillType, AntialiasType);
-
-    static void affineTo3x3(const AffineTransform&, float mat[9]);
-    static void affineTo4x4(const AffineTransform&, float mat[16]);
-    static unsigned loadShader(GraphicsContext3D*, unsigned type, const String& shaderSource);
-    static unsigned loadProgram(GraphicsContext3D*, const String& vertexShaderSource, const String& fragmentShaderSource);
-
-    GraphicsContext3D* m_context;
-    unsigned m_program;
-};
-
-}
-
-#endif // Shader_h
diff --git a/Source/WebCore/platform/graphics/gpu/SharedGraphicsContext3D.cpp b/Source/WebCore/platform/graphics/gpu/SharedGraphicsContext3D.cpp
deleted file mode 100644 (file)
index 73f3e21..0000000
+++ /dev/null
@@ -1,113 +0,0 @@
-/*
- * Copyright (C) 2011 Google Inc. All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions
- * are met:
- *
- * 1.  Redistributions of source code must retain the above copyright
- *     notice, this list of conditions and the following disclaimer.
- * 2.  Redistributions in binary form must reproduce the above copyright
- *     notice, this list of conditions and the following disclaimer in the
- *     documentation and/or other materials provided with the distribution.
- *
- * THIS SOFTWARE IS PROVIDED BY APPLE AND ITS CONTRIBUTORS "AS IS" AND ANY
- * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
- * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
- * DISCLAIMED. IN NO EVENT SHALL APPLE OR ITS CONTRIBUTORS BE LIABLE FOR ANY
- * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
- * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
- * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
- * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF
- * THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- */
-
-
-#include "config.h"
-
-#include "SharedGraphicsContext3D.h"
-
-#include "Extensions3D.h"
-#include <wtf/MainThread.h>
-
-namespace WebCore {
-
-class SharedGraphicsContext3DImpl {
-public:
-    SharedGraphicsContext3DImpl() : m_context(0) { }
-    PassRefPtr<GraphicsContext3D> getOrCreateContext()
-    {
-        bool wasCreated = false;
-
-        // If we lost the context, or can't make it current, create a new one.
-        if (m_context && (!m_context->makeContextCurrent() || (m_context->getExtensions()->getGraphicsResetStatusARB() != GraphicsContext3D::NO_ERROR)))
-            m_context.clear();
-
-        if (!m_context) {
-            createContext();
-            wasCreated = true;
-        }
-
-        if (m_context && !m_context->makeContextCurrent())
-            m_context.clear();
-
-        if (m_context && wasCreated)
-            m_context->getExtensions()->pushGroupMarkerEXT("SharedGraphicsContext");
-        return m_context;
-    }
-
-    PassRefPtr<GraphicsContext3D> getContext()
-    {
-        return m_context;
-    }
-
-    PassRefPtr<GraphicsContext3D> createContext()
-    {
-        GraphicsContext3D::Attributes attributes;
-        attributes.depth = false;
-        attributes.stencil = true;
-        attributes.antialias = false;
-        attributes.shareResources = true;
-        m_context = GraphicsContext3D::create(attributes, 0);
-        return m_context;
-    }
-private:
-    RefPtr<GraphicsContext3D> m_context;
-};
-
-PassRefPtr<GraphicsContext3D> SharedGraphicsContext3D::get()
-{
-    DEPRECATED_DEFINE_STATIC_LOCAL(SharedGraphicsContext3DImpl, impl, ());
-    return impl.getOrCreateContext();
-}
-
-enum ContextOperation {
-    Get, Create
-};
-
-static PassRefPtr<GraphicsContext3D> getOrCreateContextForImplThread(ContextOperation op)
-{
-    DEPRECATED_DEFINE_STATIC_LOCAL(SharedGraphicsContext3DImpl, impl, ());
-    return op == Create ? impl.createContext() : impl.getContext();
-}
-
-PassRefPtr<GraphicsContext3D> SharedGraphicsContext3D::getForImplThread()
-{
-    return getOrCreateContextForImplThread(Get);
-}
-
-bool SharedGraphicsContext3D::haveForImplThread()
-{
-    ASSERT(isMainThread());
-    return getOrCreateContextForImplThread(Get);
-}
-
-bool SharedGraphicsContext3D::createForImplThread()
-{
-    ASSERT(isMainThread());
-    return getOrCreateContextForImplThread(Create);
-}
-
-}
-
diff --git a/Source/WebCore/platform/graphics/gpu/SharedGraphicsContext3D.h b/Source/WebCore/platform/graphics/gpu/SharedGraphicsContext3D.h
deleted file mode 100644 (file)
index 850e63e..0000000
+++ /dev/null
@@ -1,57 +0,0 @@
-/*
- * Copyright (C) 2011 Google Inc. All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions
- * are met:
- *
- * 1.  Redistributions of source code must retain the above copyright
- *     notice, this list of conditions and the following disclaimer.
- * 2.  Redistributions in binary form must reproduce the above copyright
- *     notice, this list of conditions and the following disclaimer in the
- *     documentation and/or other materials provided with the distribution.
- *
- * THIS SOFTWARE IS PROVIDED BY APPLE AND ITS CONTRIBUTORS "AS IS" AND ANY
- * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
- * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
- * DISCLAIMED. IN NO EVENT SHALL APPLE OR ITS CONTRIBUTORS BE LIABLE FOR ANY
- * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
- * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
- * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
- * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF
- * THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- */
-
-#ifndef SharedGraphicsContext3D_h
-#define SharedGraphicsContext3D_h
-
-#include "GraphicsContext3D.h"
-#include <wtf/RefCounted.h>
-#include <wtf/RefPtr.h>
-
-namespace WebCore {
-
-class SharedGraphicsContext3D {
-public:
-    // The caller may ref this pointer, and hang onto it as long as they like.
-    // However, the context should be checked periodically to determine if it
-    // has been lost. The easiest way to do that is to simply call this
-    // function again. Note that the return value may be 0 if the
-    // GPU is unavailable.
-    static PassRefPtr<GraphicsContext3D> get();
-    // This one returns the context, and does not touch it or re-create it.
-    // Should only be called on the impl thread.
-    static PassRefPtr<GraphicsContext3D> getForImplThread();
-    // This one returns if the threaded utility context exists.
-    // Should only be called on the main thread.
-    static bool haveForImplThread();
-    // This call creates the context unconditionally, but does not touch it.
-    // Should only be called on the main thread.
-    static bool createForImplThread();
-};
-
-}
-
-#endif // SharedGraphicsContext3D_h
-
diff --git a/Source/WebCore/platform/graphics/gpu/opencl/FilterContextOpenCL.cpp b/Source/WebCore/platform/graphics/gpu/opencl/FilterContextOpenCL.cpp
deleted file mode 100644 (file)
index 6aa0503..0000000
+++ /dev/null
@@ -1,308 +0,0 @@
-/*
- * Copyright (C) 2012 University of Szeged
- * Copyright (C) 2012 Tamas Czene <tczene@inf.u-szeged.hu>
- * All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions
- * are met:
- * 1. Redistributions of source code must retain the above copyright
- *    notice, this list of conditions and the following disclaimer.
- * 2. Redistributions in binary form must reproduce the above copyright
- *    notice, this list of conditions and the following disclaimer in the
- *    documentation and/or other materials provided with the distribution.
- *
- * THIS SOFTWARE IS PROVIDED BY UNIVERSITY OF SZEGED ``AS IS'' AND ANY
- * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
- * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
- * PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL UNIVERSITY OF SZEGED OR
- * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
- * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
- * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
- * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
- * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
- * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- */
-
-#include "config.h"
-
-#if ENABLE(OPENCL)
-#include "FilterContextOpenCL.h"
-
-namespace WebCore {
-
-FilterContextOpenCL* FilterContextOpenCL::m_context = 0;
-int FilterContextOpenCL::m_alreadyInitialized = 0;
-
-FilterContextOpenCL* FilterContextOpenCL::context()
-{
-    if (m_context)
-        return m_context;
-    if (m_alreadyInitialized)
-        return 0;
-
-    m_alreadyInitialized = true;
-    FilterContextOpenCL* localContext = new FilterContextOpenCL();
-
-    // Initializing the context.
-    cl_int errorNumber;
-    cl_device_id* devices;
-    cl_platform_id firstPlatformId;
-    size_t deviceBufferSize = 0;
-
-    errorNumber = clGetPlatformIDs(1, &firstPlatformId, 0);
-    cl_context_properties contextProperties[] = { CL_CONTEXT_PLATFORM, (cl_context_properties)firstPlatformId, 0};
-    localContext->m_deviceContext = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_GPU, 0, 0, &errorNumber);
-    if (errorNumber != CL_SUCCESS) {
-        localContext->m_deviceContext = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_CPU, 0, 0, &errorNumber);
-        if (errorNumber != CL_SUCCESS)
-            return 0;
-    }
-
-    errorNumber = clGetContextInfo(localContext->m_deviceContext, CL_CONTEXT_DEVICES, 0, 0, &deviceBufferSize);
-    if (errorNumber != CL_SUCCESS)
-        return 0;
-
-    if (!deviceBufferSize)
-        return 0;
-
-    devices = reinterpret_cast<cl_device_id*>(fastMalloc(deviceBufferSize));
-    errorNumber = clGetContextInfo(localContext->m_deviceContext, CL_CONTEXT_DEVICES, deviceBufferSize, devices, 0);
-    if (errorNumber != CL_SUCCESS)
-        return 0;
-
-    localContext->m_commandQueue = clCreateCommandQueue(localContext->m_deviceContext, devices[0], 0, 0);
-    if (!localContext->m_commandQueue)
-        return 0;
-
-    localContext->m_deviceId = devices[0];
-    fastFree(devices);
-
-    cl_bool imageSupport = CL_FALSE;
-    clGetDeviceInfo(localContext->m_deviceId, CL_DEVICE_IMAGE_SUPPORT, sizeof(cl_bool), &imageSupport, 0);
-    if (imageSupport != CL_TRUE)
-        return 0;
-
-    m_context = localContext;
-    return m_context;
-}
-
-void FilterContextOpenCL::freeResources()
-{
-    clFinish(m_commandQueue);
-
-    if (m_colorMatrixWasCompiled) {
-        freeResource(m_matrixOperation);
-        freeResource(m_saturateAndHueRotateOperation);
-        freeResource(m_luminanceOperation);
-        freeResource(m_colorMatrixProgram);
-    }
-    m_colorMatrixWasCompiled = false;
-
-    if (m_turbulenceWasCompiled) {
-        freeResource(m_turbulenceOperation);
-        freeResource(m_turbulenceProgram);
-    }    
-    m_turbulenceWasCompiled = false;
-
-    if (m_transformColorSpaceWasCompiled) {
-        freeResource(m_transformColorSpaceKernel);
-        freeResource(m_transformColorSpaceProgram);
-    }
-    m_transformColorSpaceWasCompiled = false;
-}
-
-void FilterContextOpenCL::destroyContext()
-{
-    freeResources();
-
-    if (m_commandQueue)
-        clReleaseCommandQueue(m_commandQueue);
-    m_commandQueue = 0;
-
-    if (m_deviceContext)
-        clReleaseContext(m_deviceContext);
-    m_deviceContext = 0;
-
-    m_context = 0;
-}
-
-OpenCLHandle FilterContextOpenCL::createOpenCLImage(IntSize paintSize)
-{
-    FilterContextOpenCL* context = FilterContextOpenCL::context();
-
-    cl_image_format clImageFormat;
-    clImageFormat.image_channel_order = CL_RGBA;
-    clImageFormat.image_channel_data_type = CL_UNORM_INT8;
-
-#ifdef CL_API_SUFFIX__VERSION_1_2
-    cl_image_desc imageDescriptor = { CL_MEM_OBJECT_IMAGE2D, paintSize.width(), paintSize.height(), 0, 0, 0, 0, 0, 0, 0};
-    OpenCLHandle image = clCreateImage(context->deviceContext(), CL_MEM_READ_WRITE, &clImageFormat, &imageDescriptor, 0, 0);
-#else
-    OpenCLHandle image = clCreateImage2D(context->deviceContext(), CL_MEM_READ_WRITE, &clImageFormat,
-        paintSize.width(), paintSize.height(), 0, 0, 0);
-#endif
-    return image;
-}
-
-static const char* transformColorSpaceKernelProgram =
-PROGRAM(
-const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
-
-__kernel void transformColorSpace(__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);
-}
-); // End of OpenCL kernels
-
-inline bool FilterContextOpenCL::compileTransformColorSpaceProgram()
-{
-    if (m_transformColorSpaceWasCompiled || inError())
-        return !inError();
-
-    m_transformColorSpaceWasCompiled = true;
-
-    if (isResourceAllocationFailed((m_transformColorSpaceProgram = compileProgram(transformColorSpaceKernelProgram))))
-        return false;
-    if (isResourceAllocationFailed((m_transformColorSpaceKernel = kernelByName(m_transformColorSpaceProgram, "transformColorSpace"))))
-        return false;
-    return true;
-}
-
-void FilterContextOpenCL::openCLTransformColorSpace(OpenCLHandle& source, IntRect sourceSize, ColorSpace srcColorSpace, ColorSpace dstColorSpace)
-{
-    DEPRECATED_DEFINE_STATIC_LOCAL(OpenCLHandle, deviceRgbLUT, ());
-    DEPRECATED_DEFINE_STATIC_LOCAL(OpenCLHandle, linearRgbLUT, ());
-
-    if (srcColorSpace == dstColorSpace || inError())
-        return;
-
-    if ((srcColorSpace != ColorSpaceLinearRGB && srcColorSpace != ColorSpaceDeviceRGB)
-        || (dstColorSpace != ColorSpaceLinearRGB && dstColorSpace != ColorSpaceDeviceRGB))
-        return;
-
-    if (!compileTransformColorSpaceProgram())
-        return;
-
-    OpenCLHandle destination = createOpenCLImage(sourceSize.size());
-
-    RunKernel kernel(this, m_transformColorSpaceKernel, 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 = kernel.addArgument(lookUpTable.data(), sizeof(float) * 256);
-        } else
-            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 = kernel.addArgument(lookUpTable.data(), sizeof(float) * 256);
-        } else
-            kernel.addArgument(deviceRgbLUT);
-    }
-
-    kernel.run();
-    source.clear();
-    source = destination;
-}
-
-static const char* fillKernelProgram =
-PROGRAM_STR(
-__kernel void fill(__write_only image2d_t destination, float r, float g, float b, float a)
-{
-    float4 sourcePixel = (float4)(r, g, b, a);
-    write_imagef(destination, (int2)(get_global_id(0), get_global_id(1)), sourcePixel);
-}
-);
-
-inline bool FilterContextOpenCL::compileFill()
-{
-    if (m_fillWasCompiled || inError())
-        return !inError();
-
-    m_fillWasCompiled = true;
-
-    if (isResourceAllocationFailed((m_fillProgram = compileProgram(fillKernelProgram))))
-        return false;
-    if (isResourceAllocationFailed((m_fill = kernelByName(m_fillProgram, "fill"))))
-        return false;
-    return true;
-}
-
-void FilterContextOpenCL::fill(cl_mem image, IntSize imageSize, Color color)
-{
-    if (!m_context || inError())
-        return;
-
-    compileFill();
-
-    float r, g, b, a;
-
-    color.getRGBA(r, g, b, a);
-
-    RunKernel kernel(this, m_fill, imageSize.width(), imageSize.height());
-    kernel.addArgument(image);
-    kernel.addArgument(r);
-    kernel.addArgument(g);
-    kernel.addArgument(b);
-    kernel.addArgument(a);
-    kernel.run();
-}
-
-cl_program FilterContextOpenCL::compileProgram(const char* source)
-{
-    cl_program program;
-    cl_int errorNumber = 0;
-
-    program = clCreateProgramWithSource(m_deviceContext, 1, (const char**) &source, 0, &errorNumber);
-    if (isFailed(errorNumber))
-        return 0;
-
-    if (isFailed(clBuildProgram(program, 0, 0, 0, 0, 0)))
-        return 0;
-
-    return program;
-}
-
-void FilterContextOpenCL::freeResource(cl_kernel& handle)
-{
-    if (handle) {
-        clReleaseKernel(handle);
-        handle = 0;
-    }
-}
-
-void FilterContextOpenCL::freeResource(cl_program& handle)
-{
-    if (handle) {
-        clReleaseProgram(handle);
-        handle = 0;
-    }
-}
-} // namespace WebCore
-
-#endif
diff --git a/Source/WebCore/platform/graphics/gpu/opencl/FilterContextOpenCL.h b/Source/WebCore/platform/graphics/gpu/opencl/FilterContextOpenCL.h
deleted file mode 100644 (file)
index 38de758..0000000
+++ /dev/null
@@ -1,231 +0,0 @@
-/*
- * Copyright (C) 2012 University of Szeged
- * Copyright (C) 2012 Tamas Czene <tczene@inf.u-szeged.hu>
- * All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions
- * are met:
- * 1. Redistributions of source code must retain the above copyright
- *    notice, this list of conditions and the following disclaimer.
- * 2. Redistributions in binary form must reproduce the above copyright
- *    notice, this list of conditions and the following disclaimer in the
- *    documentation and/or other materials provided with the distribution.
- *
- * THIS SOFTWARE IS PROVIDED BY UNIVERSITY OF SZEGED ``AS IS'' AND ANY
- * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
- * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
- * PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL UNIVERSITY OF SZEGED OR
- * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
- * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
- * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
- * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
- * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
- * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- */
-
-#ifndef FilterContextOpenCL_h
-#define FilterContextOpenCL_h
-
-#if ENABLE(OPENCL)
-#include "CL/cl.h"
-#include "Color.h"
-#include "ColorSpace.h"
-#include "IntRect.h"
-#include "IntSize.h"
-#include "OpenCLHandle.h"
-
-#define PROGRAM_STR(...)  #__VA_ARGS__
-#define PROGRAM(...) PROGRAM_STR(__VA_ARGS__)
-
-namespace WebCore {
-
-class FilterContextOpenCL {
-public:
-    FilterContextOpenCL()
-        : m_inError(false)
-        , m_deviceId(0)
-        , m_deviceContext(0)
-        , m_commandQueue(0)
-        , m_transformColorSpaceWasCompiled(false)
-        , m_transformColorSpaceProgram(0)
-        , m_transformColorSpaceKernel(0)
-        , m_fillWasCompiled(false)
-        , m_fillProgram(0)
-        , m_fill(0)
-        , m_mergeWasCompiled(false)
-        , m_mergeProgram(0)
-        , m_mergeCopyOperation(0)
-        , m_mergeOperation(0)
-        , m_colorMatrixWasCompiled(false)
-        , m_colorMatrixProgram(0)
-        , m_matrixOperation(0)
-        , m_saturateAndHueRotateOperation(0)
-        , m_luminanceOperation(0)
-        , m_turbulenceWasCompiled(false)
-        , m_turbulenceProgram(0)
-        , m_turbulenceOperation(0)
-    {
-    }
-
-    // Returns 0 if initialization failed.
-    static FilterContextOpenCL* context();
-
-    cl_device_id deviceId() { return m_deviceId; }
-    cl_context deviceContext() { return m_deviceContext; }
-    cl_command_queue commandQueue() { return m_commandQueue; }
-
-    inline void setInError(bool errorCode = true) { m_inError = errorCode; }
-    inline bool inError() { return m_inError; }
-    inline bool isFailed(bool);
-    inline bool isResourceAllocationFailed(bool);
-
-    void freeResources();
-    void destroyContext();
-
-    OpenCLHandle createOpenCLImage(IntSize);
-
-    inline bool compileFill();
-    void fill(cl_mem, IntSize, Color);
-
-    inline bool compileTransformColorSpaceProgram();
-    void openCLTransformColorSpace(OpenCLHandle&, IntRect, ColorSpace, ColorSpace);
-
-    inline bool compileFEColorMatrix();
-    inline bool compileFETurbulence();
-    inline bool compileFEMerge();
-
-    inline void applyFEMergeCopy(OpenCLHandle, IntSize, OpenCLHandle, IntPoint&);
-    inline void applyFEMerge(OpenCLHandle, OpenCLHandle, OpenCLHandle, OpenCLHandle, IntSize, IntPoint&, IntPoint&);
-    inline void applyFEColorMatrix(OpenCLHandle, IntSize, OpenCLHandle, IntPoint, float*, int);
-    inline void applyFETurbulence(OpenCLHandle, IntSize, int, void*, void*, void*, void*, void*,
-        int*, int, int, int, int, float, float, bool, int, int);
-
-private:
-
-    class RunKernel {
-        public:
-            RunKernel(FilterContextOpenCL* context, cl_kernel kernel, size_t width, size_t height)
-                : m_context(context)
-                , m_kernel(kernel)
-                , m_index(0)
-                , m_error(context->inError())
-            {
-                m_globalSize[0] = width;
-                m_globalSize[1] = height;
-            }
-
-            void addArgument(OpenCLHandle handle)
-            {
-                if (!m_error)
-                    m_error = clSetKernelArg(m_kernel, m_index++, sizeof(OpenCLHandle), handle.handleAddress());
-            }
-
-            void addArgument(cl_int value)
-            {
-                if (!m_error)
-                    m_error = clSetKernelArg(m_kernel, m_index++, sizeof(cl_int), reinterpret_cast<void*>(&value));
-            }
-
-            void addArgument(cl_float value)
-            {
-                if (!m_error)
-                    m_error = clSetKernelArg(m_kernel, m_index++, sizeof(cl_float), reinterpret_cast<void*>(&value));
-            }
-
-            void addArgument(cl_sampler handle)
-            {
-                if (!m_error)
-                    m_error = clSetKernelArg(m_kernel, m_index++, sizeof(cl_sampler), reinterpret_cast<void*>(&handle));
-            }
-
-            OpenCLHandle addArgument(void* buffer, int size)
-            {
-                if (m_error)
-                    return 0;
-                OpenCLHandle handle(clCreateBuffer(m_context->deviceContext(), CL_MEM_READ_ONLY, size, 0, &m_error));
-                if (m_error)
-                    return 0;
-                m_error = clEnqueueWriteBuffer(m_context->commandQueue(), handle, CL_TRUE, 0, size, buffer, 0, 0, 0);
-                if (m_error)
-                    return 0;
-                m_error = clSetKernelArg(m_kernel, m_index++, sizeof(OpenCLHandle), handle.handleAddress());
-                return !m_error ? handle : 0;
-            }
-
-            void run()
-            {
-                if (m_context->isFailed(m_error))
-                    return;
-
-                m_error = clFinish(m_context->m_commandQueue);
-                if (!m_error)
-                    m_error = clEnqueueNDRangeKernel(m_context->m_commandQueue, m_kernel, 2, 0, m_globalSize, 0, 0, 0, 0);
-                m_context->isFailed(m_error);
-            }
-
-            FilterContextOpenCL* m_context;
-            cl_kernel m_kernel;
-            size_t m_globalSize[2];
-            int m_index;
-            int m_error;
-        };
-
-    cl_program compileProgram(const char*);
-    static inline cl_kernel kernelByName(cl_program program, const char* name) { return clCreateKernel(program, name, 0); }
-
-    static inline void freeResource(cl_kernel&);
-    static inline void freeResource(cl_program&);
-
-    static FilterContextOpenCL* m_context;
-    static int m_alreadyInitialized;
-    bool m_inError;
-
-    cl_device_id m_deviceId;
-    cl_context m_deviceContext;
-    cl_command_queue m_commandQueue;
-
-    bool m_transformColorSpaceWasCompiled;
-    cl_program m_transformColorSpaceProgram;
-    cl_kernel m_transformColorSpaceKernel;
-
-    bool m_fillWasCompiled;
-    cl_program m_fillProgram;
-    cl_kernel m_fill;
-
-    bool m_mergeWasCompiled;
-    cl_program m_mergeProgram;
-    cl_kernel m_mergeCopyOperation;
-    cl_kernel m_mergeOperation;
-
-    bool m_colorMatrixWasCompiled;
-    cl_program m_colorMatrixProgram;
-    cl_kernel m_matrixOperation;
-    cl_kernel m_saturateAndHueRotateOperation;
-    cl_kernel m_luminanceOperation;
-
-    bool m_turbulenceWasCompiled;
-    cl_program m_turbulenceProgram;
-    cl_kernel m_turbulenceOperation;
-};
-
-inline bool FilterContextOpenCL::isFailed(bool value)
-{
-    if (value)
-        setInError();
-    return value;
-}
-
-inline bool FilterContextOpenCL::isResourceAllocationFailed(bool value)
-{
-    if (!value)
-        setInError();
-    return !value;
-}
-
-} // namespace WebCore
-
-#endif // ENABLE(OPENCL)
-
-#endif
diff --git a/Source/WebCore/platform/graphics/gpu/opencl/OpenCLFEColorMatrix.cpp b/Source/WebCore/platform/graphics/gpu/opencl/OpenCLFEColorMatrix.cpp
deleted file mode 100644 (file)
index c6c34bd..0000000
+++ /dev/null
@@ -1,156 +0,0 @@
-/*
- * Copyright (C) 2012 University of Szeged
- * Copyright (C) 2012 Tamas Czene <tczene@inf.u-szeged.hu>
- * All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions
- * are met:
- * 1. Redistributions of source code must retain the above copyright
- *    notice, this list of conditions and the following disclaimer.
- * 2. Redistributions in binary form must reproduce the above copyright
- *    notice, this list of conditions and the following disclaimer in the
- *    documentation and/or other materials provided with the distribution.
- *
- * THIS SOFTWARE IS PROVIDED BY UNIVERSITY OF SZEGED ``AS IS'' AND ANY
- * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
- * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
- * PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL UNIVERSITY OF SZEGED OR
- * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
- * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
- * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
- * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
- * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
- * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- */
-
-#include "config.h"
-
-#if ENABLE(FILTERS) && ENABLE(OPENCL)
-#include "FEColorMatrix.h"
-
-#include "FilterContextOpenCL.h"
-
-namespace WebCore {
-
-#define COLOR_MATRIX_KERNEL(...) \
-        int2 sourceCoord = (int2) (get_global_id(0) + x, get_global_id(1) + y); \
-        float4 sourcePixel = read_imagef(source, sampler, sourceCoord); \
-        float4 destinationPixel = (float4) (__VA_ARGS__); \
-        write_imagef(destination, (int2) (get_global_id(0), get_global_id(1)), destinationPixel);
-
-static const char* colorMatrixKernelProgram =
-PROGRAM(
-const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
-
-__kernel void matrix(__write_only image2d_t destination, __read_only image2d_t source, float x, float y, __constant float *values)
-{
-    COLOR_MATRIX_KERNEL(values[0] * sourcePixel.x + values[1] * sourcePixel.y + values[2] * sourcePixel.z + values[3] * sourcePixel.w + values[4],
-        values[5] * sourcePixel.x + values[6] * sourcePixel.y + values[7] * sourcePixel.z + values[8] * sourcePixel.w + values[9],
-        values[10] * sourcePixel.x + values[11] * sourcePixel.y + values[12] * sourcePixel.z + values[13] * sourcePixel.w + values[14],
-        values[15] * sourcePixel.x + values[16] * sourcePixel.y + values[17] * sourcePixel.z + values[18] * sourcePixel.w + values[19])
-}
-
-__kernel void saturateAndHueRotate(__write_only image2d_t destination, __read_only image2d_t source, float x, float y, __constant float *components)
-{
-    COLOR_MATRIX_KERNEL(sourcePixel.x * components[0] + sourcePixel.y * components[1] + sourcePixel.z * components[2],
-        sourcePixel.x * components[3] + sourcePixel.y * components[4] + sourcePixel.z * components[5],
-        sourcePixel.x * components[6] + sourcePixel.y * components[7] + sourcePixel.z * components[8],
-        sourcePixel.w)
-}
-
-__kernel void luminance(__write_only image2d_t destination, __read_only image2d_t source, float x, float y)
-{
-    COLOR_MATRIX_KERNEL(0, 0, 0, 0.2125 * sourcePixel.x + 0.7154 * sourcePixel.y + 0.0721 * sourcePixel.z)
-}
-); // End of OpenCL kernels
-
-inline bool FilterContextOpenCL::compileFEColorMatrix()
-{
-    if (m_colorMatrixWasCompiled || inError())
-        return !inError();
-
-    m_colorMatrixWasCompiled = true;
-
-    if (isResourceAllocationFailed((m_colorMatrixProgram = compileProgram(colorMatrixKernelProgram))))
-        return false;
-    if (isResourceAllocationFailed((m_matrixOperation = kernelByName(m_colorMatrixProgram, "matrix"))))
-        return false;
-    if (isResourceAllocationFailed((m_saturateAndHueRotateOperation = kernelByName(m_colorMatrixProgram, "saturateAndHueRotate"))))
-        return false;
-    if (isResourceAllocationFailed((m_saturateAndHueRotateOperation = kernelByName(m_colorMatrixProgram, "saturateAndHueRotate"))))
-        return false;
-    if (isResourceAllocationFailed((m_luminanceOperation = kernelByName(m_colorMatrixProgram, "luminance"))))
-        return false;
-    return true;
-}
-
-inline void FilterContextOpenCL::applyFEColorMatrix(OpenCLHandle destination, IntSize destinationSize, OpenCLHandle source, IntPoint relativeSourceLocation, float* values, int type)
-{
-    cl_kernel colorMatrix;
-    OpenCLHandle clValues;
-
-    switch (type) {
-    case FECOLORMATRIX_TYPE_MATRIX:
-        colorMatrix = m_matrixOperation;
-        break;
-    case FECOLORMATRIX_TYPE_SATURATE:
-        colorMatrix = m_saturateAndHueRotateOperation;
-        break;
-    case FECOLORMATRIX_TYPE_HUEROTATE:
-        colorMatrix = m_saturateAndHueRotateOperation;
-        break;
-    case FECOLORMATRIX_TYPE_LUMINANCETOALPHA:
-        colorMatrix = m_luminanceOperation;
-        break;
-    default:
-        ASSERT_NOT_REACHED();
-        return;
-    }
-
-    RunKernel kernel(this, colorMatrix, destinationSize.width(), destinationSize.height());
-    kernel.addArgument(destination);
-    kernel.addArgument(source);
-    kernel.addArgument(relativeSourceLocation.x());
-    kernel.addArgument(relativeSourceLocation.y());
-    if (type == FECOLORMATRIX_TYPE_MATRIX)
-        clValues = kernel.addArgument(values, sizeof(float) * 20);
-    else if (type == FECOLORMATRIX_TYPE_SATURATE || type == FECOLORMATRIX_TYPE_HUEROTATE)
-        clValues = kernel.addArgument(values, sizeof(float) * 9);
-    kernel.run();
-
-    clValues.clear();
-}
-
-bool FEColorMatrix::platformApplyOpenCL()
-{
-    FilterContextOpenCL* context = FilterContextOpenCL::context();
-    if (!context)
-        return false;
-
-    if (!context->compileFEColorMatrix())
-        return true;
-
-    FilterEffect* in = inputEffect(0);
-    OpenCLHandle source = in->openCLImage();
-    OpenCLHandle destination = createOpenCLImageResult();
-
-    IntPoint relativeSourceLocation(
-        absolutePaintRect().x() - in->absolutePaintRect().location().x(),
-        absolutePaintRect().y() - in->absolutePaintRect().location().y());
-
-    float components[9];
-    if (FECOLORMATRIX_TYPE_SATURATE == m_type)
-        calculateSaturateComponents(components, m_values[0]);
-    else if (FECOLORMATRIX_TYPE_HUEROTATE == m_type)
-        calculateHueRotateComponents(components, m_values[0]);
-
-    context->applyFEColorMatrix(destination, absolutePaintRect().size(), source, relativeSourceLocation, (FECOLORMATRIX_TYPE_MATRIX == m_type) ? m_values.data() : components, m_type);
-
-    return true;
-}
-
-} // namespace WebCore
-
-#endif // ENABLE(FILTERS) && ENABLE(OPENCL)
diff --git a/Source/WebCore/platform/graphics/gpu/opencl/OpenCLFEFlood.cpp b/Source/WebCore/platform/graphics/gpu/opencl/OpenCLFEFlood.cpp
deleted file mode 100644 (file)
index 6ef1f66..0000000
+++ /dev/null
@@ -1,57 +0,0 @@
-/*
- * Copyright (C) 2013 University of Szeged
- * Copyright (C) 2013 Tamas Czene <tczene@inf.u-szeged.hu>
- * All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions
- * are met:
- * 1. Redistributions of source code must retain the above copyright
- *    notice, this list of conditions and the following disclaimer.
- * 2. Redistributions in binary form must reproduce the above copyright
- *    notice, this list of conditions and the following disclaimer in the
- *    documentation and/or other materials provided with the distribution.
- *
- * THIS SOFTWARE IS PROVIDED BY UNIVERSITY OF SZEGED ``AS IS'' AND ANY
- * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
- * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
- * PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL UNIVERSITY OF SZEGED OR
- * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
- * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
- * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
- * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
- * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
- * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- */
-
-#include "config.h"
-#if ENABLE(FILTERS) && ENABLE(OPENCL)
-#include "FEFlood.h"
-
-#include "FilterContextOpenCL.h"
-
-namespace WebCore {
-
-bool FEFlood::platformApplyOpenCL()
-{
-    FilterContextOpenCL* context = FilterContextOpenCL::context();
-
-    if (!context)
-        return false;
-
-    if (context->inError())
-        return true;
-
-    cl_mem destination = createOpenCLImageResult();
-
-    Color color = colorWithOverrideAlpha(floodColor().rgb(), floodOpacity());
-
-    context->fill(destination, absolutePaintRect().size(), color);
-
-    return true;
-}
-
-} // namespace WebCore
-
-#endif // ENABLE(FILTERS) && ENABLE(OPENCL)
diff --git a/Source/WebCore/platform/graphics/gpu/opencl/OpenCLFEImage.cpp b/Source/WebCore/platform/graphics/gpu/opencl/OpenCLFEImage.cpp
deleted file mode 100644 (file)
index 736a862..0000000
+++ /dev/null
@@ -1,51 +0,0 @@
-/*
- * Copyright (C) 2013 University of Szeged
- * Copyright (C) 2013 Tamas Czene <tczene@inf.u-szeged.hu>
- * All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions
- * are met:
- * 1. Redistributions of source code must retain the above copyright
- *    notice, this list of conditions and the following disclaimer.
- * 2. Redistributions in binary form must reproduce the above copyright
- *    notice, this list of conditions and the following disclaimer in the
- *    documentation and/or other materials provided with the distribution.
- *
- * THIS SOFTWARE IS PROVIDED BY UNIVERSITY OF SZEGED ``AS IS'' AND ANY
- * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
- * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
- * PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL UNIVERSITY OF SZEGED OR
- * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
- * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
- * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
- * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
- * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
- * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- */
-
-#include "config.h"
-#if ENABLE(FILTERS) && ENABLE(OPENCL)
-
-#include "FilterContextOpenCL.h"
-#include "SVGFEImage.h"
-#include "SVGFilter.h"
-#include "SharedBuffer.h"
-
-namespace WebCore {
-
-bool FEImage::platformApplyOpenCL()
-{
-    platformApplySoftware();
-    ImageBuffer* sourceImage = asImageBuffer();
-    if (sourceImage) {
-        RefPtr<Uint8ClampedArray> sourceImageData = sourceImage->getUnmultipliedImageData(IntRect(IntPoint(), sourceImage->internalSize()));
-        createOpenCLImageResult(sourceImageData->data());
-    }
-    return true;
-}
-
-} // namespace WebCore
-
-#endif // ENABLE(FILTERS) && ENABLE(OPENCL)
diff --git a/Source/WebCore/platform/graphics/gpu/opencl/OpenCLFEMerge.cpp b/Source/WebCore/platform/graphics/gpu/opencl/OpenCLFEMerge.cpp
deleted file mode 100644 (file)
index b2b80c2..0000000
+++ /dev/null
@@ -1,161 +0,0 @@
-/*
- * Copyright (C) 2013 University of Szeged
- * Copyright (C) 2013 Tamas Czene <tczene@inf.u-szeged.hu>
- * All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions
- * are met:
- * 1. Redistributions of source code must retain the above copyright
- *    notice, this list of conditions and the following disclaimer.
- * 2. Redistributions in binary form must reproduce the above copyright
- *    notice, this list of conditions and the following disclaimer in the
- *    documentation and/or other materials provided with the distribution.
- *
- * THIS SOFTWARE IS PROVIDED BY UNIVERSITY OF SZEGED ``AS IS'' AND ANY
- * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
- * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
- * PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL UNIVERSITY OF SZEGED OR
- * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
- * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
- * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
- * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
- * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
- * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- */
-
-#include "config.h"
-#if ENABLE(FILTERS) && ENABLE(OPENCL)
-#include "FEMerge.h"
-
-#include "FilterContextOpenCL.h"
-
-namespace WebCore {
-
-static const char* mergeKernelProgram =
-PROGRAM(
-const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
-
-__kernel void copy(__write_only image2d_t destination, __read_only image2d_t source, int x, int y)
-{
-    float4 destinationPixel = read_imagef(source, sampler, (int2) (get_global_id(0) + x, get_global_id(1) + y));
-    write_imagef(destination, (int2) (get_global_id(0), get_global_id(1)), destinationPixel);
-}
-
-__kernel void merge(__write_only image2d_t destination, __read_only image2d_t previousDestination, __read_only image2d_t sourceA, __read_only image2d_t sourceB, int xA, int yA, int xB, int yB)
-{
-    int2 destinationCoord = (int2) (get_global_id(0), get_global_id(1));
-    int2 sourceCoordA = (int2) (destinationCoord.x + xA, destinationCoord.y + yA);
-    int2 sourceCoordB = (int2) (destinationCoord.x + xB, destinationCoord.y + yB);
-    float4 destinationPixel = read_imagef(previousDestination, sampler, destinationCoord);
-    float4 sourcePixelA = read_imagef(sourceA, sampler, sourceCoordA);
-    float4 sourcePixelB = read_imagef(sourceB, sampler, sourceCoordB);
-
-    destinationPixel = sourcePixelA + destinationPixel * (1 - sourcePixelA.w);
-    destinationPixel = sourcePixelB + destinationPixel * (1 - sourcePixelB.w);
-
-    write_imagef(destination, destinationCoord, destinationPixel);
-}
-); // End of OpenCL kernels
-
-inline bool FilterContextOpenCL::compileFEMerge()
-{
-    if (m_mergeWasCompiled || inError())
-        return !inError();
-
-    m_mergeWasCompiled = true;
-
-    if (isResourceAllocationFailed((m_mergeProgram = compileProgram(mergeKernelProgram))))
-        return false;
-    if (isResourceAllocationFailed((m_mergeCopyOperation = kernelByName(m_mergeProgram, "copy"))))
-        return false;
-    if (isResourceAllocationFailed((m_mergeOperation = kernelByName(m_mergeProgram, "merge"))))
-        return false;
-    return true;
-}
-
-inline void FilterContextOpenCL::applyFEMergeCopy(OpenCLHandle destination, IntSize destinationSize, OpenCLHandle source, IntPoint& relativeSourcePoint)
-{
-    RunKernel kernel(this, m_mergeCopyOperation, destinationSize.width(), destinationSize.height());
-    kernel.addArgument(destination);
-    kernel.addArgument(source);
-    kernel.addArgument(relativeSourcePoint.x());
-    kernel.addArgument(relativeSourcePoint.y());
-    kernel.run();
-}
-
-inline void FilterContextOpenCL::applyFEMerge(OpenCLHandle destination, OpenCLHandle previousDestination, OpenCLHandle sourceA, OpenCLHandle sourceB, IntSize destinationSize, IntPoint& relativeSourcePointA, IntPoint& relativeSourcePointB)
-{
-    RunKernel kernel(this, m_mergeOperation, destinationSize.width(), destinationSize.height());
-    kernel.addArgument(destination);
-    kernel.addArgument(previousDestination);
-    kernel.addArgument(sourceA);
-    kernel.addArgument(sourceB);
-    kernel.addArgument(relativeSourcePointA.x());
-    kernel.addArgument(relativeSourcePointA.y());
-    kernel.addArgument(relativeSourcePointB.x());
-    kernel.addArgument(relativeSourcePointB.y());
-    kernel.run();
-}
-
-bool FEMerge::platformApplyOpenCL()
-{
-    FilterContextOpenCL* context = FilterContextOpenCL::context();
-    if (!context)
-        return false;
-
-    context->compileFEMerge();
-
-    unsigned size = numberOfEffectInputs();
-    ASSERT(size > 0);
-
-    OpenCLHandle destination = createOpenCLImageResult();
-    OpenCLHandle sourceA = 0;
-    OpenCLHandle sourceB = 0;
-    FilterEffect* in;
-
-    int i = 0;
-
-    if (size & 1) {
-        in = inputEffect(i++);
-        sourceA = in->openCLImage();
-        IntPoint relativeSourcePoint(in->absolutePaintRect().location());
-        relativeSourcePoint.setX(absolutePaintRect().x() - relativeSourcePoint.x());
-        relativeSourcePoint.setY(absolutePaintRect().y() - relativeSourcePoint.y());
-
-        context->applyFEMergeCopy(destination, absolutePaintRect().size(), sourceA, relativeSourcePoint);
-        if (size == 1)
-            return true;
-    } else
-        context->fill(destination, absolutePaintRect().size(), Color(0.0f, 0.0f, 0.0f, 0.0f));
-
-    OpenCLHandle previousDestination = context->createOpenCLImage(absolutePaintRect().size());
-
-    while (i < size) {
-        OpenCLHandle temp = previousDestination;
-        previousDestination = destination;
-        destination = temp;
-
-        in = inputEffect(i++);
-        sourceA = in->openCLImage();
-        IntPoint relativeSourcePointA(in->absolutePaintRect().location());
-        relativeSourcePointA.setX(absolutePaintRect().x() - relativeSourcePointA.x());
-        relativeSourcePointA.setY(absolutePaintRect().y() - relativeSourcePointA.y());
-
-        in = inputEffect(i++);
-        sourceB = in->openCLImage();
-        IntPoint relativeSourcePointB(in->absolutePaintRect().location());
-        relativeSourcePointB.setX(absolutePaintRect().x() - relativeSourcePointB.x());
-        relativeSourcePointB.setY(absolutePaintRect().y() - relativeSourcePointB.y());
-
-        context->applyFEMerge(destination, previousDestination, sourceA, sourceB, absolutePaintRect().size(), relativeSourcePointA, relativeSourcePointB);
-    }
-    setOpenCLImage(destination);
-    previousDestination.clear();
-    return true;
-}
-
-} // namespace WebCore
-
-#endif // ENABLE(FILTERS) && ENABLE(OPENCL)
diff --git a/Source/WebCore/platform/graphics/gpu/opencl/OpenCLFESourceAlpha.cpp b/Source/WebCore/platform/graphics/gpu/opencl/OpenCLFESourceAlpha.cpp
deleted file mode 100644 (file)
index ab93d07..0000000
+++ /dev/null
@@ -1,58 +0,0 @@
-/*
- * Copyright (C) 2012 University of Szeged
- * Copyright (C) 2012 Tamas Czene <tczene@inf.u-szeged.hu>
- * All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions
- * are met:
- * 1. Redistributions of source code must retain the above copyright
- *    notice, this list of conditions and the following disclaimer.
- * 2. Redistributions in binary form must reproduce the above copyright
- *    notice, this list of conditions and the following disclaimer in the
- *    documentation and/or other materials provided with the distribution.
- *
- * THIS SOFTWARE IS PROVIDED BY UNIVERSITY OF SZEGED ``AS IS'' AND ANY
- * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
- * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
- * PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL UNIVERSITY OF SZEGED OR
- * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
- * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
- * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
- * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
- * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
- * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- */
-
-#include "config.h"
-
-#if ENABLE(FILTERS) && ENABLE(OPENCL)
-#include "SourceAlpha.h"
-
-#include "Filter.h"
-#include "FilterContextOpenCL.h"
-#include "ImageBuffer.h"
-
-namespace WebCore {
-
-bool SourceAlpha::platformApplyOpenCL()
-{
-    FilterContextOpenCL* context = FilterContextOpenCL::context();
-
-    if (!context)
-        return false;
-
-    platformApplySoftware();
-    ImageBuffer* sourceImage = asImageBuffer();
-    if (!sourceImage)
-        return false;
-
-    RefPtr<Uint8ClampedArray> sourceImageData = sourceImage->getUnmultipliedImageData(IntRect(IntPoint(), sourceImage->internalSize()));
-    createOpenCLImageResult(sourceImageData->data());
-    return true;
-}
-
-} // namespace WebCore
-
-#endif // ENABLE(FILTERS) && ENABLE(OPENCL)
diff --git a/Source/WebCore/platform/graphics/gpu/opencl/OpenCLFESourceGraphic.cpp b/Source/WebCore/platform/graphics/gpu/opencl/OpenCLFESourceGraphic.cpp
deleted file mode 100644 (file)
index b2e9ef7..0000000
+++ /dev/null
@@ -1,58 +0,0 @@
-/*
- * Copyright (C) 2012 University of Szeged
- * Copyright (C) 2012 Tamas Czene <tczene@inf.u-szeged.hu>
- * All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions
- * are met:
- * 1. Redistributions of source code must retain the above copyright
- *    notice, this list of conditions and the following disclaimer.
- * 2. Redistributions in binary form must reproduce the above copyright
- *    notice, this list of conditions and the following disclaimer in the
- *    documentation and/or other materials provided with the distribution.
- *
- * THIS SOFTWARE IS PROVIDED BY UNIVERSITY OF SZEGED ``AS IS'' AND ANY
- * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
- * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
- * PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL UNIVERSITY OF SZEGED OR
- * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
- * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
- * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
- * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
- * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
- * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- */
-
-#include "config.h"
-
-#if ENABLE(FILTERS) && ENABLE(OPENCL)
-#include "SourceGraphic.h"
-
-#include "Filter.h"
-#include "FilterContextOpenCL.h"
-#include "ImageBuffer.h"
-
-namespace WebCore {
-
-bool SourceGraphic::platformApplyOpenCL()
-{
-    FilterContextOpenCL* context = FilterContextOpenCL::context();
-
-    if (!context)
-        return false;
-
-    platformApplySoftware();
-    ImageBuffer* sourceImage = asImageBuffer();
-    if (!sourceImage)
-        return false;
-
-    RefPtr<Uint8ClampedArray> sourceImageData = sourceImage->getUnmultipliedImageData(IntRect(IntPoint(), sourceImage->internalSize()));
-    createOpenCLImageResult(sourceImageData->data());
-    return true;
-}
-
-} // namespace WebCore
-
-#endif // ENABLE(FILTERS) && ENABLE(OPENCL)
diff --git a/Source/WebCore/platform/graphics/gpu/opencl/OpenCLFETurbulence.cpp b/Source/WebCore/platform/graphics/gpu/opencl/OpenCLFETurbulence.cpp
deleted file mode 100644 (file)
index d8284eb..0000000
+++ /dev/null
@@ -1,248 +0,0 @@
-/*
- * Copyright (C) 2012 University of Szeged
- * Copyright (C) 2012 Tamas Czene <tczene@inf.u-szeged.hu>
- * All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions
- * are met:
- * 1. Redistributions of source code must retain the above copyright
- *    notice, this list of conditions and the following disclaimer.
- * 2. Redistributions in binary form must reproduce the above copyright
- *    notice, this list of conditions and the following disclaimer in the
- *    documentation and/or other materials provided with the distribution.
- *
- * THIS SOFTWARE IS PROVIDED BY UNIVERSITY OF SZEGED ``AS IS'' AND ANY
- * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
- * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
- * PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL UNIVERSITY OF SZEGED OR
- * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
- * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
- * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
- * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
- * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
- * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- */
-#include "config.h"
-
-#if ENABLE(FILTERS) && ENABLE(OPENCL)
-#include "FETurbulence.h"
-
-#include "FETurbulence.cpp"
-#include "FilterContextOpenCL.h"
-#include "SVGFilter.h"
-
-namespace WebCore {
-
-static const char* turbulenceKernelProgram =
-PROGRAM(
-__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 mad(b - a, t, 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 = { 0, 0, 0, 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.
-            baseFrequencyX = (baseFrequencyX / lowFrequency < highFrequency / baseFrequencyX) ? lowFrequency : highFrequency;
-        }
-        if (baseFrequencyY) {
-            float lowFrequency = floor(tileHeight * baseFrequencyY) / tileHeight;
-            float highFrequency = ceil(tileHeight * baseFrequencyY) / tileHeight;
-            baseFrequencyY = (baseFrequencyY / lowFrequency < highFrequency / baseFrequencyY) ? lowFrequency : 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) {
-        float4 noise2DResult = (float4)( noise2D(redComponent, latticeSelector, stitchData, noiseVectorX, noiseVectorY, stitchTiles) / ratio,
-    noise2D(greenComponent, latticeSelector, stitchData, noiseVectorX, noiseVectorY, stitchTiles) / ratio,
-    noise2D(blueComponent, latticeSelector, stitchData, noiseVectorX, noiseVectorY, stitchTiles) / ratio,
-    noise2D(alphaComponent, latticeSelector, stitchData, noiseVectorX, noiseVectorY, stitchTiles) / ratio);
-
-        turbulenceFunctionResult += (type == 1) ? noise2DResult : fabs(noise2DResult);
-
-        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 = mad(0.5f, turbulenceFunctionResult, 0.5f);
-    // Clamp result.
-    turbulenceFunctionResult = clamp(turbulenceFunctionResult, 0.0f, 1.0f);
-
-    write_imagef(destination, (int2)(get_global_id(0), get_global_id(1)), turbulenceFunctionResult);
-}
-); // End of OpenCL kernels
-
-inline bool FilterContextOpenCL::compileFETurbulence()
-{
-    if (m_turbulenceWasCompiled || inError())
-        return !inError();
-
-    m_turbulenceWasCompiled = true;
-
-    if (isResourceAllocationFailed((m_turbulenceProgram = compileProgram(turbulenceKernelProgram))))
-        return false;
-    if (isResourceAllocationFailed((m_turbulenceOperation = kernelByName(m_turbulenceProgram, "Turbulence"))))
-        return false;
-    return true;
-}
-
-inline void FilterContextOpenCL::applyFETurbulence(OpenCLHandle destination,
-    IntSize destinationSize, int totalBlockSize,
-    void* transform, void* redComponent, void* greenComponent,
-    void* blueComponent, void* alphaComponent,
-    int* latticeSelector, int offsetX, int offsetY, int tileWidth, int tileHeight,
-    float baseFrequencyX, float baseFrequencyY, bool stitchTiles, int numOctaves, int type)
-{
-    RunKernel kernel(this, m_turbulenceOperation, destinationSize.width(), destinationSize.height());
-
-    kernel.addArgument(destination);
-    OpenCLHandle transformHandle(kernel.addArgument(transform, sizeof(float) * 6));
-    OpenCLHandle redComponentHandle(kernel.addArgument(redComponent, sizeof(float) * totalBlockSize * 2));
-    OpenCLHandle greenComponentHandle(kernel.addArgument(greenComponent, sizeof(float) * totalBlockSize * 2));
-    OpenCLHandle blueComponentHandle(kernel.addArgument(blueComponent, sizeof(float) * totalBlockSize * 2));
-    OpenCLHandle alphaComponentHandle(kernel.addArgument(alphaComponent, sizeof(float) * totalBlockSize * 2));
-    OpenCLHandle latticeSelectorHandle(kernel.addArgument(latticeSelector, sizeof(int) * totalBlockSize));
-    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();
-
-    transformHandle.clear();
-    redComponentHandle.clear();
-    greenComponentHandle.clear();
-    blueComponentHandle.clear();
-    alphaComponentHandle.clear();
-    latticeSelectorHandle.clear();
-}
-
-bool FETurbulence::platformApplyOpenCL()
-{
-    FilterContextOpenCL* context = FilterContextOpenCL::context();
-    if (!context)
-        return false;
-
-    if (!context->compileFETurbulence())
-        return true;
-
-    OpenCLHandle 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() };
-
-    context->applyFETurbulence(destination, absolutePaintRect().size(), 2 * s_blockSize + 2, transformComponents, paintingData.gradient,
-        paintingData.gradient + 1, paintingData.gradient + 2, paintingData.gradient + 3, paintingData.latticeSelector,
-        absolutePaintRect().x(), absolutePaintRect().y(), paintingData.filterSize.width(), paintingData.filterSize.height(),
-        m_baseFrequencyX, m_baseFrequencyY, m_stitchTiles, m_numOctaves, m_type);
-
-    return true;
-}
-
-} // namespace WebCore
-
-#endif // ENABLE(FILTERS) && ENABLE(OPENCL)
diff --git a/Source/WebCore/platform/graphics/gpu/opencl/OpenCLHandle.h b/Source/WebCore/platform/graphics/gpu/opencl/OpenCLHandle.h
deleted file mode 100644 (file)
index fa75b94..0000000
+++ /dev/null
@@ -1,65 +0,0 @@
-/*
- * Copyright (C) 2012 University of Szeged
- * Copyright (C) 2012 Tamas Czene <tczene@inf.u-szeged.hu>
- * All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions
- * are met:
- * 1. Redistributions of source code must retain the above copyright
- *    notice, this list of conditions and the following disclaimer.
- * 2. Redistributions in binary form must reproduce the above copyright
- *    notice, this list of conditions and the following disclaimer in the
- *    documentation and/or other materials provided with the distribution.
- *
- * THIS SOFTWARE IS PROVIDED BY UNIVERSITY OF SZEGED ``AS IS'' AND ANY
- * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
- * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
- * PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL UNIVERSITY OF SZEGED OR
- * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
- * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
- * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
- * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
- * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
- * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- */
-#if ENABLE(OPENCL)
-
-#ifndef OpenCLHandle_h
-#define OpenCLHandle_h
-
-#include "CL/cl.h"
-
-namespace WebCore {
-
-class OpenCLHandle {
-public:
-    OpenCLHandle() : m_openCLMemory(0) { }
-    OpenCLHandle(cl_mem openCLMemory) : m_openCLMemory(openCLMemory) { }
-
-    operator cl_mem() { return m_openCLMemory; }
-
-    void operator=(OpenCLHandle openCLMemory) { m_openCLMemory = openCLMemory; }
-
-    // This conversion operator allows implicit conversion to bool but not to other integer types.
-    typedef cl_mem (OpenCLHandle::*UnspecifiedBoolType);
-    operator UnspecifiedBoolType() const { return m_openCLMemory ? &OpenCLHandle::m_openCLMemory : 0; }
-
-    void* handleAddress() { return reinterpret_cast<void*>(&m_openCLMemory); }
-
-    void clear()
-    {
-        if (m_openCLMemory)
-            clReleaseMemObject(m_openCLMemory);
-        m_openCLMemory = 0;
-    }
-
-private:
-    cl_mem m_openCLMemory;
-};
-
-}
-
-#endif
-#endif // ENABLE(OPENCL)