Add error checking into OpenCL version of SVG filters.
authorcommit-queue@webkit.org <commit-queue@webkit.org@268f45cc-cd09-0410-ab3c-d52691b4dbfc>
Tue, 12 Feb 2013 09:40:39 +0000 (09:40 +0000)
committercommit-queue@webkit.org <commit-queue@webkit.org@268f45cc-cd09-0410-ab3c-d52691b4dbfc>
Tue, 12 Feb 2013 09:40:39 +0000 (09:40 +0000)
https://bugs.webkit.org/show_bug.cgi?id=107444

Patch by Tamas Czene <tczene@inf.u-szeged.hu> on 2013-02-12
Reviewed by Zoltan Herczeg.

In case of an error the program runs through all the remaining filters by doing nothing.
After that deletes the results of every filter and starts software rendering.

* platform/graphics/filters/FilterEffect.cpp:
(WebCore):
(WebCore::FilterEffect::applyAll): At software rendering this is a simple inline methode, but at OpenCL rendering it releases OpenCL things. If we have an error remove filter's results and start software rendering.
(WebCore::FilterEffect::clearResultsRecursive):
(WebCore::FilterEffect::openCLImageToImageBuffer):
(WebCore::FilterEffect::createOpenCLImageResult):
(WebCore::FilterEffect::transformResultColorSpace):
* platform/graphics/filters/FilterEffect.h:
(FilterEffect):
(WebCore::FilterEffect::applyAll):
* platform/graphics/gpu/opencl/FilterContextOpenCL.cpp:
(WebCore::FilterContextOpenCL::isFailed):
(WebCore):
(WebCore::FilterContextOpenCL::freeResources):
(WebCore::FilterContextOpenCL::destroyContext):
(WebCore::FilterContextOpenCL::compileTransformColorSpaceProgram):
(WebCore::FilterContextOpenCL::openCLTransformColorSpace):
(WebCore::FilterContextOpenCL::compileProgram):
(WebCore::FilterContextOpenCL::freeResource):
* platform/graphics/gpu/opencl/FilterContextOpenCL.h:
(WebCore::FilterContextOpenCL::FilterContextOpenCL):
(WebCore::FilterContextOpenCL::setInError):
(WebCore::FilterContextOpenCL::inError):
(FilterContextOpenCL):
(WebCore::FilterContextOpenCL::RunKernel::RunKernel):
(WebCore::FilterContextOpenCL::RunKernel::addArgument):
(WebCore::FilterContextOpenCL::RunKernel::run):
(RunKernel):
* platform/graphics/gpu/opencl/OpenCLFEColorMatrix.cpp:
(WebCore::FilterContextOpenCL::compileFEColorMatrix):
(WebCore::FEColorMatrix::platformApplyOpenCL):
* platform/graphics/gpu/opencl/OpenCLFETurbulence.cpp:
(WebCore::FilterContextOpenCL::compileFETurbulence):
(WebCore::FETurbulence::platformApplyOpenCL):
* rendering/svg/RenderSVGResourceFilter.cpp:
(WebCore::RenderSVGResourceFilter::postApplyResource):

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

Source/WebCore/ChangeLog
Source/WebCore/platform/graphics/filters/FilterEffect.cpp
Source/WebCore/platform/graphics/filters/FilterEffect.h
Source/WebCore/platform/graphics/gpu/opencl/FilterContextOpenCL.cpp
Source/WebCore/platform/graphics/gpu/opencl/FilterContextOpenCL.h
Source/WebCore/platform/graphics/gpu/opencl/OpenCLFEColorMatrix.cpp
Source/WebCore/platform/graphics/gpu/opencl/OpenCLFETurbulence.cpp
Source/WebCore/rendering/svg/RenderSVGResourceFilter.cpp

index 339ce6fcce2c5d1753567dccde3e859aeaa9804e..4f9c345a8909af7e98e69eab9287ab762d2ba334 100644 (file)
@@ -1,3 +1,50 @@
+2013-02-12  Tamas Czene  <tczene@inf.u-szeged.hu>
+
+        Add error checking into OpenCL version of SVG filters.
+        https://bugs.webkit.org/show_bug.cgi?id=107444
+
+        Reviewed by Zoltan Herczeg.
+
+        In case of an error the program runs through all the remaining filters by doing nothing. 
+        After that deletes the results of every filter and starts software rendering.
+
+        * platform/graphics/filters/FilterEffect.cpp:
+        (WebCore):
+        (WebCore::FilterEffect::applyAll): At software rendering this is a simple inline methode, but at OpenCL rendering it releases OpenCL things. If we have an error remove filter's results and start software rendering.
+        (WebCore::FilterEffect::clearResultsRecursive):
+        (WebCore::FilterEffect::openCLImageToImageBuffer):
+        (WebCore::FilterEffect::createOpenCLImageResult):
+        (WebCore::FilterEffect::transformResultColorSpace):
+        * platform/graphics/filters/FilterEffect.h:
+        (FilterEffect):
+        (WebCore::FilterEffect::applyAll):
+        * platform/graphics/gpu/opencl/FilterContextOpenCL.cpp:
+        (WebCore::FilterContextOpenCL::isFailed):
+        (WebCore):
+        (WebCore::FilterContextOpenCL::freeResources):
+        (WebCore::FilterContextOpenCL::destroyContext):
+        (WebCore::FilterContextOpenCL::compileTransformColorSpaceProgram):
+        (WebCore::FilterContextOpenCL::openCLTransformColorSpace):
+        (WebCore::FilterContextOpenCL::compileProgram):
+        (WebCore::FilterContextOpenCL::freeResource):
+        * platform/graphics/gpu/opencl/FilterContextOpenCL.h:
+        (WebCore::FilterContextOpenCL::FilterContextOpenCL):
+        (WebCore::FilterContextOpenCL::setInError):
+        (WebCore::FilterContextOpenCL::inError):
+        (FilterContextOpenCL):
+        (WebCore::FilterContextOpenCL::RunKernel::RunKernel):
+        (WebCore::FilterContextOpenCL::RunKernel::addArgument):
+        (WebCore::FilterContextOpenCL::RunKernel::run):
+        (RunKernel):
+        * platform/graphics/gpu/opencl/OpenCLFEColorMatrix.cpp:
+        (WebCore::FilterContextOpenCL::compileFEColorMatrix):
+        (WebCore::FEColorMatrix::platformApplyOpenCL):
+        * platform/graphics/gpu/opencl/OpenCLFETurbulence.cpp:
+        (WebCore::FilterContextOpenCL::compileFETurbulence):
+        (WebCore::FETurbulence::platformApplyOpenCL):
+        * rendering/svg/RenderSVGResourceFilter.cpp:
+        (WebCore::RenderSVGResourceFilter::postApplyResource):
+
 2013-02-12  Huang Dongsung  <luxtella@company100.net>
 
         [TexMap] Apply frames-per-second debug counter to WK1.
index 8ad8fb77637f97ed11abab241c6b762ac064a2a9..966e5347447e1544b5b5982bf0d80a5b59b86951 100644 (file)
@@ -97,6 +97,24 @@ FilterEffect* FilterEffect::inputEffect(unsigned number) const
     return m_inputEffects.at(number).get();
 }
 
+#if ENABLE(OPENCL)
+void FilterEffect::applyAll()
+{
+    if (hasResult())
+        return;
+    FilterContextOpenCL* context = FilterContextOpenCL::context();
+    if (context) {
+        apply();
+        if (!context->inError())
+            return;
+        clearResultsRecursive();
+        context->destroyContext();
+    }
+    // Software code path.
+    apply();
+}
+#endif
+
 void FilterEffect::apply()
 {
     if (hasResult())
@@ -223,6 +241,18 @@ void FilterEffect::clearResult()
 #endif
 }
 
+void FilterEffect::clearResultsRecursive()
+{
+    // Clear all results, regardless that the current effect has
+    // a result. Can be used if an effect is in an erroneous state.
+    if (hasResult())
+        clearResult();
+
+    unsigned size = m_inputEffects.size();
+    for (unsigned i = 0; i < size; ++i)
+        m_inputEffects.at(i).get()->clearResultsRecursive();
+}
+
 ImageBuffer* FilterEffect::asImageBuffer()
 {
     if (!hasResult())
@@ -248,13 +278,19 @@ ImageBuffer* FilterEffect::openCLImageToImageBuffer()
     FilterContextOpenCL* context = FilterContextOpenCL::context();
     ASSERT(context);
 
+    if (context->inError())
+        return 0;
+
     size_t origin[3] = { 0, 0, 0 };
     size_t region[3] = { m_absolutePaintRect.width(), m_absolutePaintRect.height(), 1 };
 
     RefPtr<Uint8ClampedArray> destinationPixelArray = Uint8ClampedArray::create(m_absolutePaintRect.width() * m_absolutePaintRect.height() * 4);
 
-    clFinish(context->commandQueue());
-    clEnqueueReadImage(context->commandQueue(), m_openCLImageResult, CL_TRUE, origin, region, 0, 0, destinationPixelArray->data(), 0, 0, 0);
+    if (context->isFailed(clFinish(context->commandQueue())))
+        return 0;
+
+    if (context->isFailed(clEnqueueReadImage(context->commandQueue(), m_openCLImageResult, CL_TRUE, origin, region, 0, 0, destinationPixelArray->data(), 0, 0, 0)))
+        return 0;
 
     m_imageBufferResult = ImageBuffer::create(m_absolutePaintRect.size());
     IntRect destinationRect(IntPoint(), m_absolutePaintRect.size());
@@ -426,15 +462,23 @@ Uint8ClampedArray* FilterEffect::createPremultipliedImageResult()
 #if ENABLE(OPENCL)
 OpenCLHandle FilterEffect::createOpenCLImageResult(uint8_t* source)
 {
+    FilterContextOpenCL* context = FilterContextOpenCL::context();
+    ASSERT(context);
+
+    if (context->inError())
+        return 0;
+
     ASSERT(!hasResult());
     cl_image_format clImageFormat;
     clImageFormat.image_channel_order = CL_RGBA;
     clImageFormat.image_channel_data_type = CL_UNORM_INT8;
 
-    FilterContextOpenCL* context = FilterContextOpenCL::context();
-    ASSERT(context);
+    int errorCode = 0;
     m_openCLImageResult = clCreateImage2D(context->deviceContext(), CL_MEM_READ_WRITE | (source ? CL_MEM_COPY_HOST_PTR : 0),
-        &clImageFormat, m_absolutePaintRect.width(), m_absolutePaintRect.height(), 0, source, 0);
+        &clImageFormat, m_absolutePaintRect.width(), m_absolutePaintRect.height(), 0, source, &errorCode);
+    if (context->isFailed(errorCode))
+        return 0;
+
     return m_openCLImageResult;
 }
 #endif
@@ -450,17 +494,17 @@ void FilterEffect::transformResultColorSpace(ColorSpace dstColorSpace)
 
 #if ENABLE(OPENCL)
     if (openCLImage()) {
+        if (m_imageBufferResult)
+            m_imageBufferResult.clear();
         FilterContextOpenCL* context = FilterContextOpenCL::context();
         ASSERT(context);
         context->openCLTransformColorSpace(m_openCLImageResult, absolutePaintRect(), m_resultColorSpace, dstColorSpace);
-        if (m_imageBufferResult)
-            m_imageBufferResult.clear();
     } else {
 #endif
 
-    // FIXME: We can avoid this potentially unnecessary ImageBuffer conversion by adding
-    // color space transform support for the {pre,un}multiplied arrays.
-    asImageBuffer()->transformColorSpace(m_resultColorSpace, dstColorSpace);
+        // FIXME: We can avoid this potentially unnecessary ImageBuffer conversion by adding
+        // color space transform support for the {pre,un}multiplied arrays.
+        asImageBuffer()->transformColorSpace(m_resultColorSpace, dstColorSpace);
 
 #if ENABLE(OPENCL)
     }
index 01d85a974060b8e62ebdabda0a215bc671f5d15e..f0ff887aedd6faac25a315c021ef41568a809293 100644 (file)
@@ -68,6 +68,8 @@ public:
     virtual ~FilterEffect();
 
     void clearResult();
+    void clearResultsRecursive();
+
     ImageBuffer* asImageBuffer();
     PassRefPtr<Uint8ClampedArray> asUnmultipliedImage(const IntRect&);
     PassRefPtr<Uint8ClampedArray> asPremultipliedImage(const IntRect&);
@@ -109,7 +111,12 @@ public:
     void setMaxEffectRect(const FloatRect& maxEffectRect) { m_maxEffectRect = maxEffectRect; } 
 
     void apply();
-    
+#if ENABLE(OPENCL)
+    void applyAll();
+#else
+    inline void applyAll() { apply(); }
+#endif
+
     // Correct any invalid pixels, if necessary, in the result of a filter operation.
     // This method is used to ensure valid pixel values on filter inputs and the final result.
     // Only the arithmetic composite filter ever needs to perform correction.
index 4c30596533610b404853c7e25abed355b5db09fc..580933d21dab73a569d6f36e5377048df404d88b 100644 (file)
@@ -88,6 +88,46 @@ FilterContextOpenCL* FilterContextOpenCL::context()
     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();
@@ -117,31 +157,38 @@ __kernel void transformColorSpace(__read_only image2d_t source, __write_only ima
 }
 ); // 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)
 {
     DEFINE_STATIC_LOCAL(OpenCLHandle, deviceRgbLUT, ());
     DEFINE_STATIC_LOCAL(OpenCLHandle, linearRgbLUT, ());
 
-    if (srcColorSpace == dstColorSpace)
+    if (srcColorSpace == dstColorSpace || inError())
         return;
 
     if ((srcColorSpace != ColorSpaceLinearRGB && srcColorSpace != ColorSpaceDeviceRGB)
         || (dstColorSpace != ColorSpaceLinearRGB && dstColorSpace != ColorSpaceDeviceRGB))
         return;
 
-    FilterContextOpenCL* context = FilterContextOpenCL::context();
-    ASSERT(context);
+    if (!compileTransformColorSpaceProgram())
+        return;
 
-    OpenCLHandle destination = context->createOpenCLImage(sourceSize.size());
+    OpenCLHandle destination = createOpenCLImage(sourceSize.size());
 
-    if (!m_transformColorSpaceProgram) {
-        m_transformColorSpaceProgram = compileProgram(transformColorSpaceKernelProgram);
-        ASSERT(m_transformColorSpaceProgram);
-        m_transformColorSpaceKernel = kernelByName(m_transformColorSpaceProgram, "transformColorSpace");
-        ASSERT(m_transformColorSpaceKernel);
-    }
-
-    RunKernel kernel(context, m_transformColorSpaceKernel, sourceSize.width(), sourceSize.height());
+    RunKernel kernel(this, m_transformColorSpaceKernel, sourceSize.width(), sourceSize.height());
     kernel.addArgument(source);
     kernel.addArgument(destination);
 
@@ -181,18 +228,33 @@ void FilterContextOpenCL::openCLTransformColorSpace(OpenCLHandle& source, IntRec
 cl_program FilterContextOpenCL::compileProgram(const char* source)
 {
     cl_program program;
-    cl_int errorNumber;
+    cl_int errorNumber = 0;
 
-    FilterContextOpenCL* context = FilterContextOpenCL::context();
-    ASSERT(context);
+    program = clCreateProgramWithSource(m_deviceContext, 1, (const char**) &source, 0, &errorNumber);
+    if (isFailed(errorNumber))
+        return 0;
 
-    program = clCreateProgramWithSource(context->m_deviceContext, 1, (const char**) &source, 0, 0);
-    errorNumber = clBuildProgram(program, 0, 0, 0, 0, 0);
-    if (errorNumber)
+    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
index b1898b61499570d0465711b54e3ef02854d6ac98..80264208379ae74019f6865f193f21f900134420 100644 (file)
@@ -44,17 +44,19 @@ namespace WebCore {
 class FilterContextOpenCL {
 public:
     FilterContextOpenCL()
-        : m_deviceId(0)
+        : m_inError(false)
+        , m_deviceId(0)
         , m_deviceContext(0)
         , m_commandQueue(0)
+        , m_transformColorSpaceWasCompiled(false)
         , m_transformColorSpaceProgram(0)
         , m_transformColorSpaceKernel(0)
-        , m_colorMatrixCompileStatus(openclNotCompiledYet)
+        , m_colorMatrixWasCompiled(false)
         , m_colorMatrixProgram(0)
         , m_matrixOperation(0)
         , m_saturateAndHueRotateOperation(0)
         , m_luminanceOperation(0)
-        , m_turbulenceCompileStatus(openclNotCompiledYet)
+        , m_turbulenceWasCompiled(false)
         , m_turbulenceProgram(0)
         , m_turbulenceOperation(0)
     {
@@ -67,7 +69,17 @@ public:
     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 compileTransformColorSpaceProgram();
     void openCLTransformColorSpace(OpenCLHandle&, IntRect, ColorSpace, ColorSpace);
 
     inline bool compileFEColorMatrix();
@@ -84,7 +96,8 @@ private:
             RunKernel(FilterContextOpenCL* context, cl_kernel kernel, size_t width, size_t height)
                 : m_context(context)
                 , m_kernel(kernel)
-                , index(0)
+                , m_index(0)
+                , m_error(context->inError())
             {
                 m_globalSize[0] = width;
                 m_globalSize[1] = height;
@@ -92,74 +105,103 @@ private:
 
             void addArgument(OpenCLHandle handle)
             {
-                clSetKernelArg(m_kernel, index++, sizeof(OpenCLHandle), handle.handleAddress());
+                if (!m_error)
+                    m_error = clSetKernelArg(m_kernel, m_index++, sizeof(OpenCLHandle), handle.handleAddress());
             }
 
             void addArgument(cl_int value)
             {
-                clSetKernelArg(m_kernel, index++, sizeof(cl_int), reinterpret_cast<void*>(&value));
+                if (!m_error)
+                    m_error = clSetKernelArg(m_kernel, m_index++, sizeof(cl_int), reinterpret_cast<void*>(&value));
             }
 
             void addArgument(cl_float value)
             {
-                clSetKernelArg(m_kernel, index++, sizeof(cl_float), reinterpret_cast<void*>(&value));
+                if (!m_error)
+                    m_error = clSetKernelArg(m_kernel, m_index++, sizeof(cl_float), reinterpret_cast<void*>(&value));
             }
 
             void addArgument(cl_sampler handle)
             {
-                clSetKernelArg(m_kernel, index++, sizeof(cl_sampler), reinterpret_cast<void*>(&handle));
+                if (!m_error)
+                    m_error = clSetKernelArg(m_kernel, m_index++, sizeof(cl_sampler), reinterpret_cast<void*>(&handle));
             }
 
             OpenCLHandle addArgument(void* buffer, int size)
             {
-                OpenCLHandle handle(clCreateBuffer(m_context->deviceContext(), CL_MEM_READ_ONLY, size, 0, 0));
-                clEnqueueWriteBuffer(m_context->commandQueue(), handle, CL_TRUE, 0, size, buffer, 0, 0, 0);
-                clSetKernelArg(m_kernel, index++, sizeof(OpenCLHandle), handle.handleAddress());
-                return handle;
+                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()
             {
-                clFinish(m_context->m_commandQueue);
-                clEnqueueNDRangeKernel(m_context->m_commandQueue, m_kernel, 2, 0, m_globalSize, 0, 0, 0, 0);
+                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 index;
+            int m_index;
+            int m_error;
         };
 
-    enum OpenCLCompileStatus {
-        openclNotCompiledYet,
-        openclCompileFailed,
-        openclCompileSuccessful
-    };
-
-    static cl_program compileProgram(const char*);
+    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;
 
-    OpenCLCompileStatus m_colorMatrixCompileStatus;
+    bool m_colorMatrixWasCompiled;
     cl_program m_colorMatrixProgram;
     cl_kernel m_matrixOperation;
     cl_kernel m_saturateAndHueRotateOperation; 
     cl_kernel m_luminanceOperation;
 
-    OpenCLCompileStatus m_turbulenceCompileStatus;
+    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)
index 47c88c6b1974a0c5bb2d4ccb55995865db425c7c..c6c34bdb5e4aac3b4b77d15027fc3638bd256f7d 100644 (file)
@@ -68,26 +68,22 @@ __kernel void luminance(__write_only image2d_t destination, __read_only image2d_
 
 inline bool FilterContextOpenCL::compileFEColorMatrix()
 {
-    if (m_colorMatrixCompileStatus != openclNotCompiledYet)
-        return m_colorMatrixCompileStatus == openclCompileSuccessful;
+    if (m_colorMatrixWasCompiled || inError())
+        return !inError();
 
-    m_colorMatrixCompileStatus = openclCompileFailed;
-    m_colorMatrixProgram = compileProgram(colorMatrixKernelProgram);
-    if (!m_colorMatrixProgram)
-        return false;
+    m_colorMatrixWasCompiled = true;
 
-    m_matrixOperation = kernelByName(m_colorMatrixProgram, "matrix");
-    if (!m_matrixOperation)
+    if (isResourceAllocationFailed((m_colorMatrixProgram = compileProgram(colorMatrixKernelProgram))))
         return false;
-    m_saturateAndHueRotateOperation = kernelByName(m_colorMatrixProgram, "saturateAndHueRotate");
-    if (!m_saturateAndHueRotateOperation)
+    if (isResourceAllocationFailed((m_matrixOperation = kernelByName(m_colorMatrixProgram, "matrix"))))
         return false;
-    m_luminanceOperation = kernelByName(m_colorMatrixProgram, "luminance");
-    if (!m_luminanceOperation)
+    if (isResourceAllocationFailed((m_saturateAndHueRotateOperation = kernelByName(m_colorMatrixProgram, "saturateAndHueRotate"))))
         return false;
-
-    m_colorMatrixCompileStatus = openclCompileSuccessful;
-    return openclCompileSuccessful;
+    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)
@@ -130,9 +126,12 @@ inline void FilterContextOpenCL::applyFEColorMatrix(OpenCLHandle destination, In
 bool FEColorMatrix::platformApplyOpenCL()
 {
     FilterContextOpenCL* context = FilterContextOpenCL::context();
-    if (!context || !context->compileFEColorMatrix())
+    if (!context)
         return false;
 
+    if (!context->compileFEColorMatrix())
+        return true;
+
     FilterEffect* in = inputEffect(0);
     OpenCLHandle source = in->openCLImage();
     OpenCLHandle destination = createOpenCLImageResult();
index 2c91b47bcc12251aee13e2b8dc0663747529c5b1..d8284ebc11352f545b5e1f7ab31b57a4d8e97738 100644 (file)
@@ -169,19 +169,16 @@ __kernel void Turbulence(__write_only image2d_t destination, __constant float *t
 
 inline bool FilterContextOpenCL::compileFETurbulence()
 {
-    if (m_turbulenceCompileStatus != openclNotCompiledYet)
-        return m_turbulenceCompileStatus == openclCompileSuccessful;
+    if (m_turbulenceWasCompiled || inError())
+        return !inError();
 
-    m_turbulenceCompileStatus = openclCompileFailed;
-    m_turbulenceProgram = compileProgram(turbulenceKernelProgram);
-    if (!m_turbulenceProgram)
+    m_turbulenceWasCompiled = true;
+
+    if (isResourceAllocationFailed((m_turbulenceProgram = compileProgram(turbulenceKernelProgram))))
         return false;
-    m_turbulenceOperation = kernelByName(m_turbulenceProgram, "Turbulence");
-    if (!m_turbulenceOperation)
+    if (isResourceAllocationFailed((m_turbulenceOperation = kernelByName(m_turbulenceProgram, "Turbulence"))))
         return false;
-
-    m_turbulenceCompileStatus = openclCompileSuccessful;
-    return openclCompileSuccessful;
+    return true;
 }
 
 inline void FilterContextOpenCL::applyFETurbulence(OpenCLHandle destination,
@@ -224,9 +221,12 @@ inline void FilterContextOpenCL::applyFETurbulence(OpenCLHandle destination,
 bool FETurbulence::platformApplyOpenCL()
 {
     FilterContextOpenCL* context = FilterContextOpenCL::context();
-    if (!context || !context->compileFETurbulence())
+    if (!context)
         return false;
 
+    if (!context->compileFETurbulence())
+        return true;
+
     OpenCLHandle destination = createOpenCLImageResult();
 
     PaintingData paintingData(m_seed, roundedIntSize(filterPrimitiveSubregion().size()));
index 2983c4f12572065d63f12bd32294d120347be4b1..9899096cce0002e436664264e618c4c894fb444e 100644 (file)
@@ -309,7 +309,7 @@ void RenderSVGResourceFilter::postApplyResource(RenderObject* object, GraphicsCo
         // Always true if filterData is just built (filterData->state == FilterData::Built).
         if (!lastEffect->hasResult()) {
             filterData->state = FilterData::Applying;
-            lastEffect->apply();
+            lastEffect->applyAll();
             lastEffect->correctFilterResultIfNeeded();
             lastEffect->transformResultColorSpace(ColorSpaceDeviceRGB);
         }