Changeset 142735 in webkit


Ignore:
Timestamp:
Feb 13, 2013 3:39:32 AM (11 years ago)
Author:
commit-queue@webkit.org
Message:

OpenCL implementation of Flood SVG filters.
​https://bugs.webkit.org/show_bug.cgi?id=109580

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

  • Target.pri:
  • platform/graphics/filters/FEFlood.h:

(FEFlood):

  • platform/graphics/gpu/opencl/FilterContextOpenCL.cpp:

(WebCore):
(WebCore::PROGRAM_STR):
(WebCore::FilterContextOpenCL::compileFill):
(WebCore::FilterContextOpenCL::fill):

  • platform/graphics/gpu/opencl/FilterContextOpenCL.h:

(WebCore::FilterContextOpenCL::FilterContextOpenCL):
(FilterContextOpenCL):

  • platform/graphics/gpu/opencl/OpenCLFEFlood.cpp: Added.

(WebCore):
(WebCore::FEFlood::platformApplyOpenCL):

Location:
trunk/Source/WebCore
Files:
1 added
5 edited

Legend:

Unmodified
Added
Removed

  • trunk/Source/WebCore/ChangeLog

    r142734 r142735  
     12013-02-13  Tamas Czene  <tczene@inf.u-szeged.hu>
     2
     3        OpenCL implementation of Flood SVG filters.
     4        https://bugs.webkit.org/show_bug.cgi?id=109580
     5
     6        Reviewed by Zoltan Herczeg.
     7
     8        * Target.pri:
     9        * platform/graphics/filters/FEFlood.h:
     10        (FEFlood):
     11        * platform/graphics/gpu/opencl/FilterContextOpenCL.cpp:
     12        (WebCore):
     13        (WebCore::PROGRAM_STR):
     14        (WebCore::FilterContextOpenCL::compileFill):
     15        (WebCore::FilterContextOpenCL::fill):
     16        * platform/graphics/gpu/opencl/FilterContextOpenCL.h:
     17        (WebCore::FilterContextOpenCL::FilterContextOpenCL):
     18        (FilterContextOpenCL):
     19        * platform/graphics/gpu/opencl/OpenCLFEFlood.cpp: Added.
     20        (WebCore):
     21        (WebCore::FEFlood::platformApplyOpenCL):
     22
    1232013-02-13  Mike West  <mkwst@chromium.org>
    224

  • trunk/Source/WebCore/Target.pri

    r142698 r142735  
    42204220        platform/graphics/gpu/opencl/FilterContextOpenCL.cpp \
    42214221        platform/graphics/gpu/opencl/OpenCLFEColorMatrix.cpp \
     4222        platform/graphics/gpu/opencl/OpenCLFEFlood.cpp \
    42224223        platform/graphics/gpu/opencl/OpenCLFESourceAlpha.cpp \
    42234224        platform/graphics/gpu/opencl/OpenCLFESourceGraphic.cpp \

  • trunk/Source/WebCore/platform/graphics/filters/FEFlood.h

    r97853 r142735  
    4141
    4242    virtual void platformApplySoftware();
     43#if ENABLE(OPENCL)
     44    virtual bool platformApplyOpenCL();
     45#endif
    4346    virtual void dump();
    4447

  • trunk/Source/WebCore/platform/graphics/gpu/opencl/FilterContextOpenCL.cpp

    r142596 r142735  
    226226}
    227227
     228static const char* fillKernelProgram =
     229PROGRAM_STR(
     230__kernel void fill(__write_only image2d_t destination, float r, float g, float b, float a)
     231{
     232    float4 sourcePixel = (float4)(r, g, b, a);
     233    write_imagef(destination, (int2)(get_global_id(0), get_global_id(1)), sourcePixel);
     234}
     235);
     236
     237inline bool FilterContextOpenCL::compileFill()
     238{
     239    if (m_fillWasCompiled || inError())
     240        return !inError();
     241
     242    m_fillWasCompiled = true;
     243
     244    if (isResourceAllocationFailed((m_fillProgram = compileProgram(fillKernelProgram))))
     245        return false;
     246    if (isResourceAllocationFailed((m_fill = kernelByName(m_fillProgram, "fill"))))
     247        return false;
     248    return true;
     249}
     250
     251void FilterContextOpenCL::fill(cl_mem image, IntSize imageSize, Color color)
     252{
     253    if (!m_context || inError())
     254        return;
     255
     256    compileFill();
     257
     258    float r, g, b, a;
     259
     260    color.getRGBA(r, g, b, a);
     261
     262    RunKernel kernel(this, m_fill, imageSize.width(), imageSize.height());
     263    kernel.addArgument(image);
     264    kernel.addArgument(r);
     265    kernel.addArgument(g);
     266    kernel.addArgument(b);
     267    kernel.addArgument(a);
     268    kernel.run();
     269}
     270
    228271cl_program FilterContextOpenCL::compileProgram(const char* source)
    229272{

  • trunk/Source/WebCore/platform/graphics/gpu/opencl/FilterContextOpenCL.h

    r142596 r142735  
    5252        , m_transformColorSpaceProgram(0)
    5353        , m_transformColorSpaceKernel(0)
     54        , m_fillWasCompiled(false)
     55        , m_fillProgram(0)
     56        , m_fill(0)
    5457        , m_colorMatrixWasCompiled(false)
    5558        , m_colorMatrixProgram(0)
    … …  
    8083    OpenCLHandle createOpenCLImage(IntSize);
    8184
     85    inline bool compileFill();
     86    void fill(cl_mem, IntSize, Color);
     87
    8288    inline bool compileTransformColorSpaceProgram();
    8389    void openCLTransformColorSpace(OpenCLHandle&, IntRect, ColorSpace, ColorSpace);
    … …  
    178184    cl_kernel m_transformColorSpaceKernel;
    179185
     186    bool m_fillWasCompiled;
     187    cl_program m_fillProgram;
     188    cl_kernel m_fill;
     189
    180190    bool m_colorMatrixWasCompiled;
    181191    cl_program m_colorMatrixProgram;
Note: See TracChangeset for help on using the changeset viewer.