Changeset 142596 in webkit


Ignore:
Timestamp:
Feb 12, 2013 1:40:39 AM (11 years ago)
Author:
commit-queue@webkit.org
Message:

Add error checking into OpenCL version of SVG filters.
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):

Location:
trunk/Source/WebCore
Files:
8 edited

Legend:

Unmodified
Added
Removed
  • trunk/Source/WebCore/ChangeLog

    r142595 r142596  
     12013-02-12  Tamas Czene  <tczene@inf.u-szeged.hu>
     2
     3        Add error checking into OpenCL version of SVG filters.
     4        https://bugs.webkit.org/show_bug.cgi?id=107444
     5
     6        Reviewed by Zoltan Herczeg.
     7
     8        In case of an error the program runs through all the remaining filters by doing nothing.
     9        After that deletes the results of every filter and starts software rendering.
     10
     11        * platform/graphics/filters/FilterEffect.cpp:
     12        (WebCore):
     13        (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.
     14        (WebCore::FilterEffect::clearResultsRecursive):
     15        (WebCore::FilterEffect::openCLImageToImageBuffer):
     16        (WebCore::FilterEffect::createOpenCLImageResult):
     17        (WebCore::FilterEffect::transformResultColorSpace):
     18        * platform/graphics/filters/FilterEffect.h:
     19        (FilterEffect):
     20        (WebCore::FilterEffect::applyAll):
     21        * platform/graphics/gpu/opencl/FilterContextOpenCL.cpp:
     22        (WebCore::FilterContextOpenCL::isFailed):
     23        (WebCore):
     24        (WebCore::FilterContextOpenCL::freeResources):
     25        (WebCore::FilterContextOpenCL::destroyContext):
     26        (WebCore::FilterContextOpenCL::compileTransformColorSpaceProgram):
     27        (WebCore::FilterContextOpenCL::openCLTransformColorSpace):
     28        (WebCore::FilterContextOpenCL::compileProgram):
     29        (WebCore::FilterContextOpenCL::freeResource):
     30        * platform/graphics/gpu/opencl/FilterContextOpenCL.h:
     31        (WebCore::FilterContextOpenCL::FilterContextOpenCL):
     32        (WebCore::FilterContextOpenCL::setInError):
     33        (WebCore::FilterContextOpenCL::inError):
     34        (FilterContextOpenCL):
     35        (WebCore::FilterContextOpenCL::RunKernel::RunKernel):
     36        (WebCore::FilterContextOpenCL::RunKernel::addArgument):
     37        (WebCore::FilterContextOpenCL::RunKernel::run):
     38        (RunKernel):
     39        * platform/graphics/gpu/opencl/OpenCLFEColorMatrix.cpp:
     40        (WebCore::FilterContextOpenCL::compileFEColorMatrix):
     41        (WebCore::FEColorMatrix::platformApplyOpenCL):
     42        * platform/graphics/gpu/opencl/OpenCLFETurbulence.cpp:
     43        (WebCore::FilterContextOpenCL::compileFETurbulence):
     44        (WebCore::FETurbulence::platformApplyOpenCL):
     45        * rendering/svg/RenderSVGResourceFilter.cpp:
     46        (WebCore::RenderSVGResourceFilter::postApplyResource):
     47
    1482013-02-12  Huang Dongsung  <luxtella@company100.net>
    249
  • trunk/Source/WebCore/platform/graphics/filters/FilterEffect.cpp

    r142434 r142596  
    9898}
    9999
     100#if ENABLE(OPENCL)
     101void FilterEffect::applyAll()
     102{
     103    if (hasResult())
     104        return;
     105    FilterContextOpenCL* context = FilterContextOpenCL::context();
     106    if (context) {
     107        apply();
     108        if (!context->inError())
     109            return;
     110        clearResultsRecursive();
     111        context->destroyContext();
     112    }
     113    // Software code path.
     114    apply();
     115}
     116#endif
     117
    100118void FilterEffect::apply()
    101119{
     
    224242}
    225243
     244void FilterEffect::clearResultsRecursive()
     245{
     246    // Clear all results, regardless that the current effect has
     247    // a result. Can be used if an effect is in an erroneous state.
     248    if (hasResult())
     249        clearResult();
     250
     251    unsigned size = m_inputEffects.size();
     252    for (unsigned i = 0; i < size; ++i)
     253        m_inputEffects.at(i).get()->clearResultsRecursive();
     254}
     255
    226256ImageBuffer* FilterEffect::asImageBuffer()
    227257{
     
    249279    ASSERT(context);
    250280
     281    if (context->inError())
     282        return 0;
     283
    251284    size_t origin[3] = { 0, 0, 0 };
    252285    size_t region[3] = { m_absolutePaintRect.width(), m_absolutePaintRect.height(), 1 };
     
    254287    RefPtr<Uint8ClampedArray> destinationPixelArray = Uint8ClampedArray::create(m_absolutePaintRect.width() * m_absolutePaintRect.height() * 4);
    255288
    256     clFinish(context->commandQueue());
    257     clEnqueueReadImage(context->commandQueue(), m_openCLImageResult, CL_TRUE, origin, region, 0, 0, destinationPixelArray->data(), 0, 0, 0);
     289    if (context->isFailed(clFinish(context->commandQueue())))
     290        return 0;
     291
     292    if (context->isFailed(clEnqueueReadImage(context->commandQueue(), m_openCLImageResult, CL_TRUE, origin, region, 0, 0, destinationPixelArray->data(), 0, 0, 0)))
     293        return 0;
    258294
    259295    m_imageBufferResult = ImageBuffer::create(m_absolutePaintRect.size());
     
    427463OpenCLHandle FilterEffect::createOpenCLImageResult(uint8_t* source)
    428464{
     465    FilterContextOpenCL* context = FilterContextOpenCL::context();
     466    ASSERT(context);
     467
     468    if (context->inError())
     469        return 0;
     470
    429471    ASSERT(!hasResult());
    430472    cl_image_format clImageFormat;
     
    432474    clImageFormat.image_channel_data_type = CL_UNORM_INT8;
    433475
    434     FilterContextOpenCL* context = FilterContextOpenCL::context();
    435     ASSERT(context);
     476    int errorCode = 0;
    436477    m_openCLImageResult = clCreateImage2D(context->deviceContext(), CL_MEM_READ_WRITE | (source ? CL_MEM_COPY_HOST_PTR : 0),
    437         &clImageFormat, m_absolutePaintRect.width(), m_absolutePaintRect.height(), 0, source, 0);
     478        &clImageFormat, m_absolutePaintRect.width(), m_absolutePaintRect.height(), 0, source, &errorCode);
     479    if (context->isFailed(errorCode))
     480        return 0;
     481
    438482    return m_openCLImageResult;
    439483}
     
    451495#if ENABLE(OPENCL)
    452496    if (openCLImage()) {
     497        if (m_imageBufferResult)
     498            m_imageBufferResult.clear();
    453499        FilterContextOpenCL* context = FilterContextOpenCL::context();
    454500        ASSERT(context);
    455501        context->openCLTransformColorSpace(m_openCLImageResult, absolutePaintRect(), m_resultColorSpace, dstColorSpace);
    456         if (m_imageBufferResult)
    457             m_imageBufferResult.clear();
    458502    } else {
    459503#endif
    460504
    461     // FIXME: We can avoid this potentially unnecessary ImageBuffer conversion by adding
    462     // color space transform support for the {pre,un}multiplied arrays.
    463     asImageBuffer()->transformColorSpace(m_resultColorSpace, dstColorSpace);
     505        // FIXME: We can avoid this potentially unnecessary ImageBuffer conversion by adding
     506        // color space transform support for the {pre,un}multiplied arrays.
     507        asImageBuffer()->transformColorSpace(m_resultColorSpace, dstColorSpace);
    464508
    465509#if ENABLE(OPENCL)
  • trunk/Source/WebCore/platform/graphics/filters/FilterEffect.h

    r138250 r142596  
    6969
    7070    void clearResult();
     71    void clearResultsRecursive();
     72
    7173    ImageBuffer* asImageBuffer();
    7274    PassRefPtr<Uint8ClampedArray> asUnmultipliedImage(const IntRect&);
     
    110112
    111113    void apply();
    112    
     114#if ENABLE(OPENCL)
     115    void applyAll();
     116#else
     117    inline void applyAll() { apply(); }
     118#endif
     119
    113120    // Correct any invalid pixels, if necessary, in the result of a filter operation.
    114121    // This method is used to ensure valid pixel values on filter inputs and the final result.
  • trunk/Source/WebCore/platform/graphics/gpu/opencl/FilterContextOpenCL.cpp

    r137999 r142596  
    8989}
    9090
     91void FilterContextOpenCL::freeResources()
     92{
     93    clFinish(m_commandQueue);
     94
     95    if (m_colorMatrixWasCompiled) {
     96        freeResource(m_matrixOperation);
     97        freeResource(m_saturateAndHueRotateOperation);
     98        freeResource(m_luminanceOperation);
     99        freeResource(m_colorMatrixProgram);
     100    }
     101    m_colorMatrixWasCompiled = false;
     102
     103    if (m_turbulenceWasCompiled) {
     104        freeResource(m_turbulenceOperation);
     105        freeResource(m_turbulenceProgram);
     106    }   
     107    m_turbulenceWasCompiled = false;
     108
     109    if (m_transformColorSpaceWasCompiled) {
     110        freeResource(m_transformColorSpaceKernel);
     111        freeResource(m_transformColorSpaceProgram);
     112    }
     113    m_transformColorSpaceWasCompiled = false;
     114}
     115
     116void FilterContextOpenCL::destroyContext()
     117{
     118    freeResources();
     119
     120    if (m_commandQueue)
     121        clReleaseCommandQueue(m_commandQueue);
     122    m_commandQueue = 0;
     123
     124    if (m_deviceContext)
     125        clReleaseContext(m_deviceContext);
     126    m_deviceContext = 0;
     127
     128    m_context = 0;
     129}
     130
    91131OpenCLHandle FilterContextOpenCL::createOpenCLImage(IntSize paintSize)
    92132{
     
    118158); // End of OpenCL kernels
    119159
     160inline bool FilterContextOpenCL::compileTransformColorSpaceProgram()
     161{
     162    if (m_transformColorSpaceWasCompiled || inError())
     163        return !inError();
     164
     165    m_transformColorSpaceWasCompiled = true;
     166
     167    if (isResourceAllocationFailed((m_transformColorSpaceProgram = compileProgram(transformColorSpaceKernelProgram))))
     168        return false;
     169    if (isResourceAllocationFailed((m_transformColorSpaceKernel = kernelByName(m_transformColorSpaceProgram, "transformColorSpace"))))
     170        return false;
     171    return true;
     172}
     173
    120174void FilterContextOpenCL::openCLTransformColorSpace(OpenCLHandle& source, IntRect sourceSize, ColorSpace srcColorSpace, ColorSpace dstColorSpace)
    121175{
     
    123177    DEFINE_STATIC_LOCAL(OpenCLHandle, linearRgbLUT, ());
    124178
    125     if (srcColorSpace == dstColorSpace)
     179    if (srcColorSpace == dstColorSpace || inError())
    126180        return;
    127181
     
    130184        return;
    131185
    132     FilterContextOpenCL* context = FilterContextOpenCL::context();
    133     ASSERT(context);
    134 
    135     OpenCLHandle destination = context->createOpenCLImage(sourceSize.size());
    136 
    137     if (!m_transformColorSpaceProgram) {
    138         m_transformColorSpaceProgram = compileProgram(transformColorSpaceKernelProgram);
    139         ASSERT(m_transformColorSpaceProgram);
    140         m_transformColorSpaceKernel = kernelByName(m_transformColorSpaceProgram, "transformColorSpace");
    141         ASSERT(m_transformColorSpaceKernel);
    142     }
    143 
    144     RunKernel kernel(context, m_transformColorSpaceKernel, sourceSize.width(), sourceSize.height());
     186    if (!compileTransformColorSpaceProgram())
     187        return;
     188
     189    OpenCLHandle destination = createOpenCLImage(sourceSize.size());
     190
     191    RunKernel kernel(this, m_transformColorSpaceKernel, sourceSize.width(), sourceSize.height());
    145192    kernel.addArgument(source);
    146193    kernel.addArgument(destination);
     
    182229{
    183230    cl_program program;
    184     cl_int errorNumber;
    185 
    186     FilterContextOpenCL* context = FilterContextOpenCL::context();
    187     ASSERT(context);
    188 
    189     program = clCreateProgramWithSource(context->m_deviceContext, 1, (const char**) &source, 0, 0);
    190     errorNumber = clBuildProgram(program, 0, 0, 0, 0, 0);
    191     if (errorNumber)
     231    cl_int errorNumber = 0;
     232
     233    program = clCreateProgramWithSource(m_deviceContext, 1, (const char**) &source, 0, &errorNumber);
     234    if (isFailed(errorNumber))
     235        return 0;
     236
     237    if (isFailed(clBuildProgram(program, 0, 0, 0, 0, 0)))
    192238        return 0;
    193239
    194240    return program;
    195241}
     242
     243void FilterContextOpenCL::freeResource(cl_kernel& handle)
     244{
     245    if (handle) {
     246        clReleaseKernel(handle);
     247        handle = 0;
     248    }
     249}
     250
     251void FilterContextOpenCL::freeResource(cl_program& handle)
     252{
     253    if (handle) {
     254        clReleaseProgram(handle);
     255        handle = 0;
     256    }
     257}
    196258} // namespace WebCore
    197259
  • trunk/Source/WebCore/platform/graphics/gpu/opencl/FilterContextOpenCL.h

    r137999 r142596  
    4545public:
    4646    FilterContextOpenCL()
    47         : m_deviceId(0)
     47        : m_inError(false)
     48        , m_deviceId(0)
    4849        , m_deviceContext(0)
    4950        , m_commandQueue(0)
     51        , m_transformColorSpaceWasCompiled(false)
    5052        , m_transformColorSpaceProgram(0)
    5153        , m_transformColorSpaceKernel(0)
    52         , m_colorMatrixCompileStatus(openclNotCompiledYet)
     54        , m_colorMatrixWasCompiled(false)
    5355        , m_colorMatrixProgram(0)
    5456        , m_matrixOperation(0)
    5557        , m_saturateAndHueRotateOperation(0)
    5658        , m_luminanceOperation(0)
    57         , m_turbulenceCompileStatus(openclNotCompiledYet)
     59        , m_turbulenceWasCompiled(false)
    5860        , m_turbulenceProgram(0)
    5961        , m_turbulenceOperation(0)
     
    6870    cl_command_queue commandQueue() { return m_commandQueue; }
    6971
     72    inline void setInError(bool errorCode = true) { m_inError = errorCode; }
     73    inline bool inError() { return m_inError; }
     74    inline bool isFailed(bool);
     75    inline bool isResourceAllocationFailed(bool);
     76
     77    void freeResources();
     78    void destroyContext();
     79
    7080    OpenCLHandle createOpenCLImage(IntSize);
     81
     82    inline bool compileTransformColorSpaceProgram();
    7183    void openCLTransformColorSpace(OpenCLHandle&, IntRect, ColorSpace, ColorSpace);
    7284
     
    8597                : m_context(context)
    8698                , m_kernel(kernel)
    87                 , index(0)
     99                , m_index(0)
     100                , m_error(context->inError())
    88101            {
    89102                m_globalSize[0] = width;
     
    93106            void addArgument(OpenCLHandle handle)
    94107            {
    95                 clSetKernelArg(m_kernel, index++, sizeof(OpenCLHandle), handle.handleAddress());
     108                if (!m_error)
     109                    m_error = clSetKernelArg(m_kernel, m_index++, sizeof(OpenCLHandle), handle.handleAddress());
    96110            }
    97111
    98112            void addArgument(cl_int value)
    99113            {
    100                 clSetKernelArg(m_kernel, index++, sizeof(cl_int), reinterpret_cast<void*>(&value));
     114                if (!m_error)
     115                    m_error = clSetKernelArg(m_kernel, m_index++, sizeof(cl_int), reinterpret_cast<void*>(&value));
    101116            }
    102117
    103118            void addArgument(cl_float value)
    104119            {
    105                 clSetKernelArg(m_kernel, index++, sizeof(cl_float), reinterpret_cast<void*>(&value));
     120                if (!m_error)
     121                    m_error = clSetKernelArg(m_kernel, m_index++, sizeof(cl_float), reinterpret_cast<void*>(&value));
    106122            }
    107123
    108124            void addArgument(cl_sampler handle)
    109125            {
    110                 clSetKernelArg(m_kernel, index++, sizeof(cl_sampler), reinterpret_cast<void*>(&handle));
     126                if (!m_error)
     127                    m_error = clSetKernelArg(m_kernel, m_index++, sizeof(cl_sampler), reinterpret_cast<void*>(&handle));
    111128            }
    112129
    113130            OpenCLHandle addArgument(void* buffer, int size)
    114131            {
    115                 OpenCLHandle handle(clCreateBuffer(m_context->deviceContext(), CL_MEM_READ_ONLY, size, 0, 0));
    116                 clEnqueueWriteBuffer(m_context->commandQueue(), handle, CL_TRUE, 0, size, buffer, 0, 0, 0);
    117                 clSetKernelArg(m_kernel, index++, sizeof(OpenCLHandle), handle.handleAddress());
    118                 return handle;
     132                if (m_error)
     133                    return 0;
     134                OpenCLHandle handle(clCreateBuffer(m_context->deviceContext(), CL_MEM_READ_ONLY, size, 0, &m_error));
     135                if (m_error)
     136                    return 0;
     137                m_error = clEnqueueWriteBuffer(m_context->commandQueue(), handle, CL_TRUE, 0, size, buffer, 0, 0, 0);
     138                if (m_error)
     139                    return 0;
     140                m_error = clSetKernelArg(m_kernel, m_index++, sizeof(OpenCLHandle), handle.handleAddress());
     141                return !m_error ? handle : 0;
    119142            }
    120143
    121144            void run()
    122145            {
    123                 clFinish(m_context->m_commandQueue);
    124                 clEnqueueNDRangeKernel(m_context->m_commandQueue, m_kernel, 2, 0, m_globalSize, 0, 0, 0, 0);
     146                if (m_context->isFailed(m_error))
     147                    return;
     148
     149                m_error = clFinish(m_context->m_commandQueue);
     150                if (!m_error)
     151                    m_error = clEnqueueNDRangeKernel(m_context->m_commandQueue, m_kernel, 2, 0, m_globalSize, 0, 0, 0, 0);
     152                m_context->isFailed(m_error);
    125153            }
    126154
     
    128156            cl_kernel m_kernel;
    129157            size_t m_globalSize[2];
    130             int index;
     158            int m_index;
     159            int m_error;
    131160        };
    132161
    133     enum OpenCLCompileStatus {
    134         openclNotCompiledYet,
    135         openclCompileFailed,
    136         openclCompileSuccessful
    137     };
    138 
    139     static cl_program compileProgram(const char*);
     162    cl_program compileProgram(const char*);
    140163    static inline cl_kernel kernelByName(cl_program program, const char* name) { return clCreateKernel(program, name, 0); }
     164
     165    static inline void freeResource(cl_kernel&);
     166    static inline void freeResource(cl_program&);
    141167
    142168    static FilterContextOpenCL* m_context;
    143169    static int m_alreadyInitialized;
     170    bool m_inError;
    144171
    145172    cl_device_id m_deviceId;
     
    147174    cl_command_queue m_commandQueue;
    148175
     176    bool m_transformColorSpaceWasCompiled;
    149177    cl_program m_transformColorSpaceProgram;
    150178    cl_kernel m_transformColorSpaceKernel;
    151179
    152     OpenCLCompileStatus m_colorMatrixCompileStatus;
     180    bool m_colorMatrixWasCompiled;
    153181    cl_program m_colorMatrixProgram;
    154182    cl_kernel m_matrixOperation;
     
    156184    cl_kernel m_luminanceOperation;
    157185
    158     OpenCLCompileStatus m_turbulenceCompileStatus;
     186    bool m_turbulenceWasCompiled;
    159187    cl_program m_turbulenceProgram;
    160188    cl_kernel m_turbulenceOperation;
    161189};
    162190
     191inline bool FilterContextOpenCL::isFailed(bool value)
     192{
     193    if (value)
     194        setInError();
     195    return value;
     196}
     197
     198inline bool FilterContextOpenCL::isResourceAllocationFailed(bool value)
     199{
     200    if (!value)
     201        setInError();
     202    return !value;
     203}
     204
    163205} // namespace WebCore
    164206
  • trunk/Source/WebCore/platform/graphics/gpu/opencl/OpenCLFEColorMatrix.cpp

    r137999 r142596  
    6969inline bool FilterContextOpenCL::compileFEColorMatrix()
    7070{
    71     if (m_colorMatrixCompileStatus != openclNotCompiledYet)
    72         return m_colorMatrixCompileStatus == openclCompileSuccessful;
     71    if (m_colorMatrixWasCompiled || inError())
     72        return !inError();
    7373
    74     m_colorMatrixCompileStatus = openclCompileFailed;
    75     m_colorMatrixProgram = compileProgram(colorMatrixKernelProgram);
    76     if (!m_colorMatrixProgram)
     74    m_colorMatrixWasCompiled = true;
     75
     76    if (isResourceAllocationFailed((m_colorMatrixProgram = compileProgram(colorMatrixKernelProgram))))
    7777        return false;
    78 
    79     m_matrixOperation = kernelByName(m_colorMatrixProgram, "matrix");
    80     if (!m_matrixOperation)
     78    if (isResourceAllocationFailed((m_matrixOperation = kernelByName(m_colorMatrixProgram, "matrix"))))
    8179        return false;
    82     m_saturateAndHueRotateOperation = kernelByName(m_colorMatrixProgram, "saturateAndHueRotate");
    83     if (!m_saturateAndHueRotateOperation)
     80    if (isResourceAllocationFailed((m_saturateAndHueRotateOperation = kernelByName(m_colorMatrixProgram, "saturateAndHueRotate"))))
    8481        return false;
    85     m_luminanceOperation = kernelByName(m_colorMatrixProgram, "luminance");
    86     if (!m_luminanceOperation)
     82    if (isResourceAllocationFailed((m_saturateAndHueRotateOperation = kernelByName(m_colorMatrixProgram, "saturateAndHueRotate"))))
    8783        return false;
    88 
    89     m_colorMatrixCompileStatus = openclCompileSuccessful;
    90     return openclCompileSuccessful;
     84    if (isResourceAllocationFailed((m_luminanceOperation = kernelByName(m_colorMatrixProgram, "luminance"))))
     85        return false;
     86    return true;
    9187}
    9288
     
    131127{
    132128    FilterContextOpenCL* context = FilterContextOpenCL::context();
    133     if (!context || !context->compileFEColorMatrix())
     129    if (!context)
    134130        return false;
     131
     132    if (!context->compileFEColorMatrix())
     133        return true;
    135134
    136135    FilterEffect* in = inputEffect(0);
  • trunk/Source/WebCore/platform/graphics/gpu/opencl/OpenCLFETurbulence.cpp

    r137999 r142596  
    170170inline bool FilterContextOpenCL::compileFETurbulence()
    171171{
    172     if (m_turbulenceCompileStatus != openclNotCompiledYet)
    173         return m_turbulenceCompileStatus == openclCompileSuccessful;
    174 
    175     m_turbulenceCompileStatus = openclCompileFailed;
    176     m_turbulenceProgram = compileProgram(turbulenceKernelProgram);
    177     if (!m_turbulenceProgram)
     172    if (m_turbulenceWasCompiled || inError())
     173        return !inError();
     174
     175    m_turbulenceWasCompiled = true;
     176
     177    if (isResourceAllocationFailed((m_turbulenceProgram = compileProgram(turbulenceKernelProgram))))
    178178        return false;
    179     m_turbulenceOperation = kernelByName(m_turbulenceProgram, "Turbulence");
    180     if (!m_turbulenceOperation)
     179    if (isResourceAllocationFailed((m_turbulenceOperation = kernelByName(m_turbulenceProgram, "Turbulence"))))
    181180        return false;
    182 
    183     m_turbulenceCompileStatus = openclCompileSuccessful;
    184     return openclCompileSuccessful;
     181    return true;
    185182}
    186183
     
    225222{
    226223    FilterContextOpenCL* context = FilterContextOpenCL::context();
    227     if (!context || !context->compileFETurbulence())
     224    if (!context)
    228225        return false;
     226
     227    if (!context->compileFETurbulence())
     228        return true;
    229229
    230230    OpenCLHandle destination = createOpenCLImageResult();
  • trunk/Source/WebCore/rendering/svg/RenderSVGResourceFilter.cpp

    r138835 r142596  
    310310        if (!lastEffect->hasResult()) {
    311311            filterData->state = FilterData::Applying;
    312             lastEffect->apply();
     312            lastEffect->applyAll();
    313313            lastEffect->correctFilterResultIfNeeded();
    314314            lastEffect->transformResultColorSpace(ColorSpaceDeviceRGB);
Note: See TracChangeset for help on using the changeset viewer.