wg_engine: Introduce fill effect for webgpu renderer

Issue: https://github.com/thorvg/thorvg/issues/3054

src/renderer/wg_engine/tvgWgCompositor.cpp

@@ -757,7 +757,7 @@ bool WgCompositor::gaussianBlur(WgContext& context, WgRenderStorage* dst, const
     assert(params->rd);
     assert(compose->rdViewport);
     assert(!renderPassEncoder);
-    auto renderDataGaussian = (WgRenderDataGaussian*)params->rd;
+    auto renderData = (WgRenderDataEffectParams*)params->rd;
     auto aabb = compose->aabb;
     auto viewport = compose->rdViewport;
     WgRenderStorage* sbuff = dst;
@@ -766,12 +766,12 @@ bool WgCompositor::gaussianBlur(WgContext& context, WgRenderStorage* dst, const
     // begin compute pass
     WGPUComputePassDescriptor computePassDesc{ .label = "Compute pass gaussian blur" };
     WGPUComputePassEncoder computePassEncoder = wgpuCommandEncoderBeginComputePass(commandEncoder, &computePassDesc);
-    for (uint32_t level = 0; level < renderDataGaussian->level; level++) {
+    for (uint32_t level = 0; level < renderData->level; level++) {
         // horizontal blur
         if (params->direction != 2) {
             wgpuComputePassEncoderSetBindGroup(computePassEncoder, 0, sbuff->bindGroupRead, 0, nullptr);
             wgpuComputePassEncoderSetBindGroup(computePassEncoder, 1, dbuff->bindGroupWrite, 0, nullptr);
-            wgpuComputePassEncoderSetBindGroup(computePassEncoder, 2, renderDataGaussian->bindGroupGaussian, 0, nullptr);
+            wgpuComputePassEncoderSetBindGroup(computePassEncoder, 2, renderData->bindGroupParams, 0, nullptr);
             wgpuComputePassEncoderSetBindGroup(computePassEncoder, 3, viewport->bindGroupViewport, 0, nullptr);
             wgpuComputePassEncoderSetPipeline(computePassEncoder, pipelines.gaussian_horz);
             wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, (aabb.w - 1) / 128 + 1, aabb.h, 1);
@@ -781,7 +781,7 @@ bool WgCompositor::gaussianBlur(WgContext& context, WgRenderStorage* dst, const
         if (params->direction != 1) {
             wgpuComputePassEncoderSetBindGroup(computePassEncoder, 0, sbuff->bindGroupRead, 0, nullptr);
             wgpuComputePassEncoderSetBindGroup(computePassEncoder, 1, dbuff->bindGroupWrite, 0, nullptr);
-            wgpuComputePassEncoderSetBindGroup(computePassEncoder, 2, renderDataGaussian->bindGroupGaussian, 0, nullptr);
+            wgpuComputePassEncoderSetBindGroup(computePassEncoder, 2, renderData->bindGroupParams, 0, nullptr);
             wgpuComputePassEncoderSetBindGroup(computePassEncoder, 3, viewport->bindGroupViewport, 0, nullptr);
             wgpuComputePassEncoderSetPipeline(computePassEncoder, pipelines.gaussian_vert);
             wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, aabb.w, (aabb.h - 1) / 128 + 1, 1);
@@ -808,7 +808,7 @@ bool WgCompositor::dropShadow(WgContext& context, WgRenderStorage* dst, const Re
     assert(compose->rdViewport);
     assert(!renderPassEncoder);
 
-    auto renderDataDropShadow = (WgRenderDataDropShadow*)params->rd;
+    auto renderDataParams = (WgRenderDataEffectParams*)params->rd;
     auto aabb = compose->aabb;
     auto viewport = compose->rdViewport;
 
@@ -821,7 +821,7 @@ bool WgCompositor::dropShadow(WgContext& context, WgRenderStorage* dst, const Re
         // horizontal blur
         wgpuComputePassEncoderSetBindGroup(computePassEncoder, 0, sbuff->bindGroupRead, 0, nullptr);
         wgpuComputePassEncoderSetBindGroup(computePassEncoder, 1, dbuff->bindGroupWrite, 0, nullptr);
-        wgpuComputePassEncoderSetBindGroup(computePassEncoder, 2, renderDataDropShadow->bindGroupGaussian, 0, nullptr);
+        wgpuComputePassEncoderSetBindGroup(computePassEncoder, 2, renderDataParams->bindGroupParams, 0, nullptr);
         wgpuComputePassEncoderSetBindGroup(computePassEncoder, 3, viewport->bindGroupViewport, 0, nullptr);
         wgpuComputePassEncoderSetPipeline(computePassEncoder, pipelines.gaussian_horz);
         wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, (aabb.w - 1) / 128 + 1, aabb.h, 1);
@@ -829,7 +829,7 @@ bool WgCompositor::dropShadow(WgContext& context, WgRenderStorage* dst, const Re
         // vertical blur
         wgpuComputePassEncoderSetBindGroup(computePassEncoder, 0, sbuff->bindGroupRead, 0, nullptr);
         wgpuComputePassEncoderSetBindGroup(computePassEncoder, 1, dbuff->bindGroupWrite, 0, nullptr);
-        wgpuComputePassEncoderSetBindGroup(computePassEncoder, 2, renderDataDropShadow->bindGroupGaussian, 0, nullptr);
+        wgpuComputePassEncoderSetBindGroup(computePassEncoder, 2, renderDataParams->bindGroupParams, 0, nullptr);
         wgpuComputePassEncoderSetBindGroup(computePassEncoder, 3, viewport->bindGroupViewport, 0, nullptr);
         wgpuComputePassEncoderSetPipeline(computePassEncoder, pipelines.gaussian_vert);
         wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, aabb.w, (aabb.h - 1) / 128 + 1, 1);
@@ -844,7 +844,7 @@ bool WgCompositor::dropShadow(WgContext& context, WgRenderStorage* dst, const Re
         WGPUComputePassEncoder computePassEncoder = wgpuCommandEncoderBeginComputePass(commandEncoder, &computePassDesc);
         wgpuComputePassEncoderSetBindGroup(computePassEncoder, 0, bindGroupStorageTemp, 0, nullptr);
         wgpuComputePassEncoderSetBindGroup(computePassEncoder, 1, dst->bindGroupWrite, 0, nullptr);
-        wgpuComputePassEncoderSetBindGroup(computePassEncoder, 2, renderDataDropShadow->bindGroupDropShadow, 0, nullptr);
+        wgpuComputePassEncoderSetBindGroup(computePassEncoder, 2, renderDataParams->bindGroupParams, 0, nullptr);
         wgpuComputePassEncoderSetBindGroup(computePassEncoder, 3, viewport->bindGroupViewport, 0, nullptr);
         wgpuComputePassEncoderSetPipeline(computePassEncoder, pipelines.dropshadow);
         wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, (aabb.w - 1) / 128 + 1, aabb.h, 1);
@@ -854,3 +854,31 @@ bool WgCompositor::dropShadow(WgContext& context, WgRenderStorage* dst, const Re
 
     return true;
 }
+
+
+bool WgCompositor::fillEffect(WgContext& context, WgRenderStorage* dst, const RenderEffectFill* params, const WgCompose* compose)
+{
+    assert(dst);
+    assert(params);
+    assert(params->rd);
+    assert(compose->rdViewport);
+    assert(!renderPassEncoder);
+
+    auto renderDataParams = (WgRenderDataEffectParams*)params->rd;
+    auto aabb = compose->aabb;
+    auto viewport = compose->rdViewport;
+
+    copyTexture(&storageTemp0, dst, aabb);
+    WGPUComputePassDescriptor computePassDesc{ .label = "Compute pass fill" };
+    WGPUComputePassEncoder computePassEncoder = wgpuCommandEncoderBeginComputePass(commandEncoder, &computePassDesc);
+    wgpuComputePassEncoderSetBindGroup(computePassEncoder, 0, bindGroupStorageTemp, 0, nullptr);
+    wgpuComputePassEncoderSetBindGroup(computePassEncoder, 1, dst->bindGroupWrite, 0, nullptr);
+    wgpuComputePassEncoderSetBindGroup(computePassEncoder, 2, renderDataParams->bindGroupParams, 0, nullptr);
+    wgpuComputePassEncoderSetBindGroup(computePassEncoder, 3, viewport->bindGroupViewport, 0, nullptr);
+    wgpuComputePassEncoderSetPipeline(computePassEncoder, pipelines.fill_effect);
+    wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, (aabb.w - 1) / 128 + 1, aabb.h, 1);
+    wgpuComputePassEncoderEnd(computePassEncoder);
+    wgpuComputePassEncoderRelease(computePassEncoder);
+
+    return true;
+}

src/renderer/wg_engine/tvgWgCompositor.h

@@ -114,6 +114,7 @@ public:
     // effects
     bool gaussianBlur(WgContext& context, WgRenderStorage* dst, const RenderEffectGaussianBlur* params, const WgCompose* compose);
     bool dropShadow(WgContext& context, WgRenderStorage* dst, const RenderEffectDropShadow* params, const WgCompose* compose);
+    bool fillEffect(WgContext& context, WgRenderStorage* dst, const RenderEffectFill* params, const WgCompose* compose);
 };
 
 #endif // _TVG_WG_COMPOSITOR_H_

src/renderer/wg_engine/tvgWgPipelines.cpp

@@ -184,8 +184,8 @@ void WgPipelines::initialize(WgContext& context)
     // bind group layouts blit
     const WGPUBindGroupLayout bindGroupLayoutsBlit[] { layouts.layoutTexSampled };
     // bind group layouts effects
-    const WGPUBindGroupLayout bindGroupLayoutsGaussian[] { layouts.layoutTexStrorage1RO, layouts.layoutTexStrorage1WO, layouts.layoutBuffer1Un, layouts.layoutBuffer1Un };
+    const WGPUBindGroupLayout bindGroupLayoutsGauss[] { layouts.layoutTexStrorage1RO, layouts.layoutTexStrorage1WO, layouts.layoutBuffer1Un, layouts.layoutBuffer1Un };
-    const WGPUBindGroupLayout bindGroupLayoutsDropShadow[] { layouts.layoutTexStrorage2RO, layouts.layoutTexStrorage1WO, layouts.layoutBuffer1Un, layouts.layoutBuffer1Un };
+    const WGPUBindGroupLayout bindGroupLayoutsEffects[] { layouts.layoutTexStrorage2RO, layouts.layoutTexStrorage1WO, layouts.layoutBuffer1Un, layouts.layoutBuffer1Un };
 
     // depth stencil state markup
     const WGPUDepthStencilState depthStencilStateNonZero = makeDepthStencilState(WGPUCompareFunction_Always, false, WGPUCompareFunction_Always, WGPUStencilOperation_IncrementWrap, WGPUCompareFunction_Always, WGPUStencilOperation_DecrementWrap);
@@ -222,8 +222,8 @@ void WgPipelines::initialize(WgContext& context)
     // shader blit
     shader_blit = createShaderModule(context.device, "The shader blit", cShaderSrc_Blit);
     // shader effects
-    shader_gaussian = createShaderModule(context.device, "The shader gaussian", cShaderSrc_GaussianBlur);
+    shader_gauss = createShaderModule(context.device, "The shader effects", cShaderSrc_GaussianBlur);
-    shader_dropshadow = createShaderModule(context.device, "The shader drop shadow", cShaderSrc_DropShadow);
+    shader_effects = createShaderModule(context.device, "The shader effects", cShaderSrc_Effects);
 
     // layouts
     layout_stencil = createPipelineLayout(context.device, bindGroupLayoutsStencil, 2);
@@ -243,8 +243,8 @@ void WgPipelines::initialize(WgContext& context)
     // layout blit
     layout_blit = createPipelineLayout(context.device, bindGroupLayoutsBlit, 1);
     // layout effects
-    layout_gaussian = createPipelineLayout(context.device, bindGroupLayoutsGaussian, 4);
+    layout_gauss = createPipelineLayout(context.device, bindGroupLayoutsGauss, 4);
-    layout_dropshadow = createPipelineLayout(context.device, bindGroupLayoutsDropShadow, 4);
+    layout_effects = createPipelineLayout(context.device, bindGroupLayoutsEffects, 4);
 
     // render pipeline nonzero
     nonzero = createRenderPipeline(
@@ -450,15 +450,16 @@ void WgPipelines::initialize(WgContext& context)
         depthStencilStateScene, multisampleStateX1);
 
     // compute pipeline gaussian blur
-    gaussian_horz = createComputePipeline(context.device, "The compute pipeline gaussian blur horizontal", shader_gaussian, "cs_main_horz", layout_gaussian);
+    gaussian_horz = createComputePipeline(context.device, "The compute pipeline gaussian blur horizontal", shader_gauss, "cs_main_horz", layout_gauss);
-    gaussian_vert = createComputePipeline(context.device, "The compute pipeline gaussian blur vertical",   shader_gaussian, "cs_main_vert", layout_gaussian);
+    gaussian_vert = createComputePipeline(context.device, "The compute pipeline gaussian blur vertical",   shader_gauss, "cs_main_vert", layout_gauss);
-    // compute pipeline drop shadow
+    dropshadow    = createComputePipeline(context.device, "The compute pipeline drop shadow blend", shader_effects, "cs_main_drop_shadow", layout_effects);
-    dropshadow = createComputePipeline(context.device, "The compute pipeline drop shadow blend", shader_dropshadow, "cs_main", layout_dropshadow);
+    fill_effect   = createComputePipeline(context.device, "The compute pipeline fill effect", shader_effects, "cs_main_fill", layout_effects);
 }
 
 void WgPipelines::releaseGraphicHandles(WgContext& context)
 {
     // pipeline effects
+    releaseComputePipeline(fill_effect);
     releaseComputePipeline(dropshadow);
     releaseComputePipeline(gaussian_vert);
     releaseComputePipeline(gaussian_horz);
@@ -492,8 +493,8 @@ void WgPipelines::releaseGraphicHandles(WgContext& context)
     releaseRenderPipeline(evenodd);
     releaseRenderPipeline(nonzero);
     // layouts
-    releasePipelineLayout(layout_dropshadow);
+    releasePipelineLayout(layout_effects);
-    releasePipelineLayout(layout_gaussian);
+    releasePipelineLayout(layout_gauss);
     releasePipelineLayout(layout_blit);
     releasePipelineLayout(layout_scene_compose);
     releasePipelineLayout(layout_scene_blend);
@@ -507,8 +508,7 @@ void WgPipelines::releaseGraphicHandles(WgContext& context)
     releasePipelineLayout(layout_depth);
     releasePipelineLayout(layout_stencil);
     // shaders
-    releaseShaderModule(shader_dropshadow);
+    releaseShaderModule(shader_effects);
-    releaseShaderModule(shader_gaussian);
     releaseShaderModule(shader_blit);
     releaseShaderModule(shader_scene_compose);
     releaseShaderModule(shader_scene_blend);

src/renderer/wg_engine/tvgWgPipelines.h

@@ -47,8 +47,8 @@ private:
     // shader blit
     WGPUShaderModule shader_blit{};
     // shader effects
-    WGPUShaderModule shader_gaussian;
+    WGPUShaderModule shader_gauss;
-    WGPUShaderModule shader_dropshadow;
+    WGPUShaderModule shader_effects;
 
     // layouts helpers
     WGPUPipelineLayout layout_stencil{};
@@ -68,8 +68,8 @@ private:
     // layouts blit
     WGPUPipelineLayout layout_blit{};
     // layouts effects
-    WGPUPipelineLayout layout_gaussian{};
+    WGPUPipelineLayout layout_gauss{};
-    WGPUPipelineLayout layout_dropshadow{};
+    WGPUPipelineLayout layout_effects{};
 public:
     // pipelines stencil markup
     WGPURenderPipeline nonzero{};
@@ -101,6 +101,7 @@ public:
     WGPUComputePipeline gaussian_horz{};
     WGPUComputePipeline gaussian_vert{};
     WGPUComputePipeline dropshadow{};
+    WGPUComputePipeline fill_effect{}; 
 private:
     void releaseGraphicHandles(WgContext& context);
     WGPUShaderModule createShaderModule(WGPUDevice device, const char* label, const char* code);

src/renderer/wg_engine/tvgWgRenderData.cpp

@@ -580,133 +580,81 @@ void WgRenderDataViewportPool::release(WgContext& context)
 }
 
 //***********************************************************************
-// WgRenderDataGaussian
+// WgRenderDataEffectParams
 //***********************************************************************
 
-void WgRenderDataGaussian::update(WgContext& context, RenderEffectGaussianBlur* gaussian, const Matrix& transform)
+void WgRenderDataEffectParams::update(WgContext& context, const WgShaderTypeEffectParams& effectParams)
+{
+    if (context.allocateBufferUniform(bufferParams, &effectParams.params, sizeof(effectParams.params))) {
+        context.layouts.releaseBindGroup(bindGroupParams);
+        bindGroupParams = context.layouts.createBindGroupBuffer1Un(bufferParams);
+    }
+}
+
+
+void WgRenderDataEffectParams::update(WgContext& context, const RenderEffectGaussianBlur* gaussian, const Matrix& transform)
 {
     assert(gaussian);
-    // compute gaussian blur data
+    WgShaderTypeEffectParams effectParams;
-    WgShaderTypeGaussianBlur gaussianSettings;
+    effectParams.update(gaussian, transform);
-    gaussianSettings.update(gaussian, transform);
+    update(context, effectParams);
-    // update bind group and buffers
-    bool bufferSettingsChanged = context.allocateBufferUniform(bufferSettings, &gaussianSettings.settings, sizeof(gaussianSettings.settings));
-    if (bufferSettingsChanged) {
-        // update bind group
-        context.layouts.releaseBindGroup(bindGroupGaussian);
-        bindGroupGaussian = context.layouts.createBindGroupBuffer1Un(bufferSettings);
-    }
     level = int(WG_GAUSSIAN_MAX_LEVEL * ((gaussian->quality - 1) * 0.01f)) + 1;
-    extend = gaussianSettings.extend;
+    extend = effectParams.extend;
 }
 
 
-void WgRenderDataGaussian::release(WgContext& context)
+void WgRenderDataEffectParams::update(WgContext& context, const RenderEffectDropShadow* dropShadow, const Matrix& transform)
-{
-    context.releaseBuffer(bufferSettings);
-    context.layouts.releaseBindGroup(bindGroupGaussian);
-}
-
-//***********************************************************************
-// WgRenderDataGaussianPool
-//***********************************************************************
-
-WgRenderDataGaussian* WgRenderDataGaussianPool::allocate(WgContext& context)
-{
-    WgRenderDataGaussian* renderData{};
-    if (mPool.count > 0) {
-        renderData = mPool.last();
-        mPool.pop();
-    } else {
-        renderData = new WgRenderDataGaussian();
-        mList.push(renderData);
-    }
-    return renderData;
-}
-
-
-void WgRenderDataGaussianPool::free(WgContext& context, WgRenderDataGaussian* renderData)
-{
-    if (renderData) mPool.push(renderData);
-}
-
-
-void WgRenderDataGaussianPool::release(WgContext& context)
-{
-    ARRAY_FOREACH(p, mList) {
-        (*p)->release(context);
-        delete(*p);
-    }
-    mPool.clear();
-    mList.clear();
-}
-
-//***********************************************************************
-// WgRenderDataDropShadow
-//***********************************************************************
-
-void WgRenderDataDropShadow::update(WgContext& context, RenderEffectDropShadow* dropShadow, const Matrix& transform)
 {
     assert(dropShadow);
-    // compute gaussian blur data
+    WgShaderTypeEffectParams effectParams;
-    WgShaderTypeGaussianBlur gaussianSettings;
+    effectParams.update(dropShadow, transform);
-    gaussianSettings.update(dropShadow, transform);
+    update(context, effectParams);
-    // update bind group and buffers
-    bool bufferGaussianChanged = context.allocateBufferUniform(bufferGaussian, &gaussianSettings.settings, sizeof(gaussianSettings.settings));
-    if (bufferGaussianChanged) {
-        // update bind group
-        context.layouts.releaseBindGroup(bindGroupGaussian);
-        bindGroupGaussian = context.layouts.createBindGroupBuffer1Un(bufferGaussian);
-    }
-    // compute drop shadow data
-    WgShaderTypeDropShadow dropShadowSettings;
-    dropShadowSettings.update(dropShadow, transform);
-    // update bind group and buffers
-    bool bufferSettingsChanged = context.allocateBufferUniform(bufferSettings, &dropShadowSettings.settings, sizeof(dropShadowSettings.settings));
-    if (bufferSettingsChanged) {
-        // update bind group
-        context.layouts.releaseBindGroup(bindGroupDropShadow);
-        bindGroupDropShadow = context.layouts.createBindGroupBuffer1Un(bufferSettings);
-    }
     level = int(WG_GAUSSIAN_MAX_LEVEL * ((dropShadow->quality - 1) * 0.01f)) + 1;
-    extend = gaussianSettings.extend;
+    extend = effectParams.extend;
-    offset = dropShadowSettings.offset;
+    offset = effectParams.offset;
 }
 
 
-void WgRenderDataDropShadow::release(WgContext& context)
+void WgRenderDataEffectParams::update(WgContext& context, const RenderEffectFill* fill)
 {
-    context.releaseBuffer(bufferSettings);
+    assert(fill);
-    context.releaseBuffer(bufferGaussian);
+    WgShaderTypeEffectParams effectParams;
-    context.layouts.releaseBindGroup(bindGroupDropShadow);
+    effectParams.update(fill);
-    context.layouts.releaseBindGroup(bindGroupGaussian);
+    update(context, effectParams);
+}
+
+
+void WgRenderDataEffectParams::release(WgContext& context)
+{
+    context.releaseBuffer(bufferParams);
+    context.layouts.releaseBindGroup(bindGroupParams);
 }
 
 //***********************************************************************
-// WgRenderDataGaussianPool
+// WgRenderDataColorsPool
 //***********************************************************************
 
-WgRenderDataDropShadow* WgRenderDataDropShadowPool::allocate(WgContext& context)
+WgRenderDataEffectParams* WgRenderDataEffectParamsPool::allocate(WgContext& context)
 {
-    WgRenderDataDropShadow* renderData{};
+    WgRenderDataEffectParams* renderData{};
     if (mPool.count > 0) {
         renderData = mPool.last();
         mPool.pop();
     } else {
-        renderData = new WgRenderDataDropShadow();
+        renderData = new WgRenderDataEffectParams();
         mList.push(renderData);
     }
     return renderData;
 }
 
 
-void WgRenderDataDropShadowPool::free(WgContext& context, WgRenderDataDropShadow* renderData)
+void WgRenderDataEffectParamsPool::free(WgContext& context, WgRenderDataEffectParams* renderData)
 {
     if (renderData) mPool.push(renderData);
 }
 
 
-void WgRenderDataDropShadowPool::release(WgContext& context)
+void WgRenderDataEffectParamsPool::release(WgContext& context)
 {
     ARRAY_FOREACH(p, mList) {
         (*p)->release(context);

src/renderer/wg_engine/tvgWgRenderData.h

@@ -25,6 +25,7 @@
 
 #include "tvgWgPipelines.h"
 #include "tvgWgGeometry.h"
+#include "tvgWgShaderTypes.h"
 
 struct WgMeshData {
     WGPUBuffer bufferPosition{};
@@ -190,55 +191,34 @@ public:
     void release(WgContext& context);
 };
 
+// gaussian blur, drop shadow, fill, tint, tritone
 #define WG_GAUSSIAN_MAX_LEVEL 3
-struct WgRenderDataGaussian
+struct WgRenderDataEffectParams
 {
-    WGPUBindGroup bindGroupGaussian{};
+    WGPUBindGroup bindGroupParams{};
-    WGPUBuffer bufferSettings{};
+    WGPUBuffer bufferParams{};
-    uint32_t extend{};
-    uint32_t level{};
-
-    void update(WgContext& context, RenderEffectGaussianBlur* gaussian, const Matrix& transform);
-    void release(WgContext& context);
-};
-
-class WgRenderDataGaussianPool {
-private:
-    // pool contains all created but unused render data for gaussian filter
-    Array<WgRenderDataGaussian*> mPool;
-    // list contains all created render data for gaussian filter
-    // to ensure that all created instances will be released
-    Array<WgRenderDataGaussian*> mList;
-public:
-    WgRenderDataGaussian* allocate(WgContext& context);
-    void free(WgContext& context, WgRenderDataGaussian* renderData);
-    void release(WgContext& context);
-};
-
-struct WgRenderDataDropShadow
-{
-    WGPUBindGroup bindGroupGaussian{};
-    WGPUBindGroup bindGroupDropShadow{};
-    WGPUBuffer bufferGaussian{};
-    WGPUBuffer bufferSettings{};
     uint32_t extend{};
     uint32_t level{};
     Point offset{};
 
-    void update(WgContext& context, RenderEffectDropShadow* dropShadow, const Matrix& transform);
+    void update(WgContext& context, const WgShaderTypeEffectParams& effectParams);
+    void update(WgContext& context, const RenderEffectGaussianBlur* gaussian, const Matrix& transform);
+    void update(WgContext& context, const RenderEffectDropShadow* dropShadow, const Matrix& transform);
+    void update(WgContext& context, const RenderEffectFill* fill);
     void release(WgContext& context);
 };
 
-class WgRenderDataDropShadowPool {
+// effect params pool
+class WgRenderDataEffectParamsPool {
 private:
-    // pool contains all created but unused render data for drop shadow
+    // pool contains all created but unused render data for params
-    Array<WgRenderDataDropShadow*> mPool;
+    Array<WgRenderDataEffectParams*> mPool;
-    // list contains all created render data for drop shadow
+    // list contains all created render data for params
     // to ensure that all created instances will be released
-    Array<WgRenderDataDropShadow*> mList;
+    Array<WgRenderDataEffectParams*> mList;
 public:
-    WgRenderDataDropShadow* allocate(WgContext& context);
+    WgRenderDataEffectParams* allocate(WgContext& context);
-    void free(WgContext& context, WgRenderDataDropShadow* renderData);
+    void free(WgContext& context, WgRenderDataEffectParams* renderData);
     void release(WgContext& context);
 };
 

src/renderer/wg_engine/tvgWgRenderer.cpp

@@ -52,9 +52,8 @@ void WgRenderer::release()
     // clear render data paint pools
     mRenderDataShapePool.release(mContext);
     mRenderDataPicturePool.release(mContext);
-    mRenderDataGaussianPool.release(mContext);
-    mRenderDataDropShadowPool.release(mContext);
     mRenderDataViewportPool.release(mContext);
+    mRenderDataEffectParamsPool.release(mContext);
     WgMeshDataPool::gMeshDataPool->release(mContext);
 
     // clear render storage pool
@@ -524,24 +523,35 @@ void WgRenderer::prepare(RenderEffect* effect, const Matrix& transform)
 {
     // prepare gaussian blur data
     if (effect->type == SceneEffect::GaussianBlur) {
-        auto gaussianBlur = (RenderEffectGaussianBlur*)effect;
+        auto renderEffect = (RenderEffectGaussianBlur*)effect;
-        auto renderDataGaussian = (WgRenderDataGaussian*)gaussianBlur->rd;
+        auto renderData = (WgRenderDataEffectParams*)renderEffect->rd;
-        if (!renderDataGaussian) {
+        if (!renderData) {
-            renderDataGaussian = mRenderDataGaussianPool.allocate(mContext);
+            renderData = mRenderDataEffectParamsPool.allocate(mContext);
-            gaussianBlur->rd = renderDataGaussian;
+            renderEffect->rd = renderData;
         }
-        renderDataGaussian->update(mContext, gaussianBlur, transform);
+        renderData->update(mContext, renderEffect, transform);
         effect->valid = true;
     } else
     // prepare drop shadow data
     if (effect->type == SceneEffect::DropShadow) {
-        auto dropShadow = (RenderEffectDropShadow*)effect;
+        auto renderEffect = (RenderEffectDropShadow*)effect;
-        auto renderDataDropShadow = (WgRenderDataDropShadow*)dropShadow->rd;
+        auto renderData = (WgRenderDataEffectParams*)renderEffect->rd;
-        if (!renderDataDropShadow) {
+        if (!renderData) {
-            renderDataDropShadow = mRenderDataDropShadowPool.allocate(mContext);
+            renderData = mRenderDataEffectParamsPool.allocate(mContext);
-            dropShadow->rd = renderDataDropShadow;
+            renderEffect->rd = renderData;
         }
-        renderDataDropShadow->update(mContext, dropShadow, transform);
+        renderData->update(mContext, renderEffect, transform);
+        effect->valid = true;
+    }
+    // prepare fill
+    if (effect->type == SceneEffect::Fill) {
+        auto renderEffect = (RenderEffectFill*)effect;
+        auto renderData = (WgRenderDataEffectParams*)renderEffect->rd;
+        if (!renderData) {
+            renderData = mRenderDataEffectParamsPool.allocate(mContext);
+            renderEffect->rd = renderData;
+        }
+        renderData->update(mContext, renderEffect);
         effect->valid = true;
     }
 }
@@ -552,25 +562,25 @@ bool WgRenderer::region(RenderEffect* effect)
     // update gaussian blur region
     if (effect->type == SceneEffect::GaussianBlur) {
         auto gaussian = (RenderEffectGaussianBlur*)effect;
-        auto renderDataGaussian = (WgRenderDataGaussian*)gaussian->rd;
+        auto renderData = (WgRenderDataEffectParams*)gaussian->rd;
         if (gaussian->direction != 2) {
-            gaussian->extend.x = -renderDataGaussian->extend;
+            gaussian->extend.x = -renderData->extend;
-            gaussian->extend.w = +renderDataGaussian->extend * 2;
+            gaussian->extend.w = +renderData->extend * 2;
         }
         if (gaussian->direction != 1) {
-            gaussian->extend.y = -renderDataGaussian->extend;
+            gaussian->extend.y = -renderData->extend;
-            gaussian->extend.h = +renderDataGaussian->extend * 2;
+            gaussian->extend.h = +renderData->extend * 2;
         }
         return true;
     } else
     // update drop shadow region
     if (effect->type == SceneEffect::DropShadow) {
         auto dropShadow = (RenderEffectDropShadow*)effect;
-        auto renderDataDropShadow = (WgRenderDataDropShadow*)dropShadow->rd;
+        auto renderData = (WgRenderDataEffectParams*)dropShadow->rd;
-        dropShadow->extend.x = -(renderDataDropShadow->extend + std::abs(renderDataDropShadow->offset.x));
+        dropShadow->extend.x = -(renderData->extend + std::abs(renderData->offset.x));
-        dropShadow->extend.w = +(renderDataDropShadow->extend + std::abs(renderDataDropShadow->offset.x)) * 2;
+        dropShadow->extend.w = +(renderData->extend + std::abs(renderData->offset.x)) * 2;
-        dropShadow->extend.y = -(renderDataDropShadow->extend + std::abs(renderDataDropShadow->offset.y));
+        dropShadow->extend.y = -(renderData->extend + std::abs(renderData->offset.y));
-        dropShadow->extend.h = +(renderDataDropShadow->extend + std::abs(renderDataDropShadow->offset.y)) * 2;
+        dropShadow->extend.h = +(renderData->extend + std::abs(renderData->offset.y)) * 2;
         return true;
     }
     return false;
@@ -587,6 +597,7 @@ bool WgRenderer::render(RenderCompositor* cmp, const RenderEffect* effect, TVG_U
     switch (effect->type) {
         case SceneEffect::GaussianBlur: return mCompositor.gaussianBlur(mContext, dst, (RenderEffectGaussianBlur*)effect, comp);
         case SceneEffect::DropShadow: return mCompositor.dropShadow(mContext, dst, (RenderEffectDropShadow*)effect, comp);
+        case SceneEffect::Fill: return mCompositor.fillEffect(mContext, dst, (RenderEffectFill*)effect, comp);
         default: return false;
     }
     return false;
@@ -595,11 +606,8 @@ bool WgRenderer::render(RenderCompositor* cmp, const RenderEffect* effect, TVG_U
 
 void WgRenderer::dispose(RenderEffect* effect)
 {
-    switch (effect->type) {
+    auto renderData = (WgRenderDataEffectParams*)effect->rd;
-        case SceneEffect::GaussianBlur: mRenderDataGaussianPool.free(mContext, (WgRenderDataGaussian*)effect->rd);
+    mRenderDataEffectParamsPool.free(mContext, renderData);
-        case SceneEffect::DropShadow: mRenderDataDropShadowPool.free(mContext, (WgRenderDataDropShadow*)effect->rd);
-        default: effect->rd = nullptr;
-    }
     effect->rd = nullptr;
 };
 

src/renderer/wg_engine/tvgWgRenderer.h

@@ -85,9 +85,8 @@ private:
     // render data paint pools
     WgRenderDataShapePool mRenderDataShapePool;
     WgRenderDataPicturePool mRenderDataPicturePool;
-    WgRenderDataGaussianPool mRenderDataGaussianPool;
-    WgRenderDataDropShadowPool mRenderDataDropShadowPool;
     WgRenderDataViewportPool mRenderDataViewportPool;
+    WgRenderDataEffectParamsPool mRenderDataEffectParamsPool;
 
     // rendering context
     WgContext mContext;

src/renderer/wg_engine/tvgWgShaderSrc.cpp

@@ -711,7 +711,7 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f {
 const char* cShaderSrc_GaussianBlur = R"(
 @group(0) @binding(0) var imageSrc : texture_storage_2d<rgba8unorm, read>;
 @group(1) @binding(0) var imageDst : texture_storage_2d<rgba8unorm, write>;
-@group(2) @binding(0) var<uniform> settings: vec4f;
+@group(2) @binding(0) var<uniform> settings: array<vec4f, 3>;
 @group(3) @binding(0) var<uniform> viewport: vec4f;
 
 const N: u32 = 128;
@@ -728,9 +728,9 @@ fn gaussian(x: f32, sigma: f32) -> f32 {
 fn cs_main_horz(@builtin(global_invocation_id) gid: vec3u,
                 @builtin(local_invocation_id)  lid: vec3u) {
     // settings decode
-    let sigma = settings.x;
+    let sigma = settings[0].x;
-    let scale = settings.y;
+    let scale = settings[0].y;
-    let size = i32(settings.z);
+    let size = i32(settings[0].z);
 
     // viewport decode
     let xmin = i32(viewport.x);
@@ -771,9 +771,9 @@ fn cs_main_horz(@builtin(global_invocation_id) gid: vec3u,
 fn cs_main_vert(@builtin(global_invocation_id) gid: vec3u,
                 @builtin(local_invocation_id)  lid: vec3u) {
     // settings decode
-    let sigma = settings.x;
+    let sigma = settings[0].x;
-    let scale = settings.y;
+    let scale = settings[0].y;
-    let size = i32(settings.z);
+    let size = i32(settings[0].z);
 
     // viewport decode
     let xmin = i32(viewport.x);
@@ -811,19 +811,19 @@ fn cs_main_vert(@builtin(global_invocation_id) gid: vec3u,
 }
 )";
 
-const char* cShaderSrc_DropShadow = R"(
+const char* cShaderSrc_Effects = R"(
 @group(0) @binding(0) var imageSrc : texture_storage_2d<rgba8unorm, read>;
 @group(0) @binding(1) var imageSdw : texture_storage_2d<rgba8unorm, read>;
 @group(1) @binding(0) var imageTrg : texture_storage_2d<rgba8unorm, write>;
-@group(2) @binding(0) var<uniform> settings: array<vec4f, 2>;
+@group(2) @binding(0) var<uniform> settings: array<vec4f, 3>;
 @group(3) @binding(0) var<uniform> viewport: vec4f;
 
 @compute @workgroup_size(128, 1)
-fn cs_main(@builtin(global_invocation_id) gid: vec3u) {
+fn cs_main_drop_shadow(@builtin(global_invocation_id) gid: vec3u) {
     // decode viewport and settings
     let vmin = vec2u(viewport.xy);
     let vmax = vec2u(viewport.zw);
-    let voff = vec2i(settings[1].xy);
+    let voff = vec2i(settings[2].xy);
 
     // tex coord
     let uid = gid.xy + vmin;
@@ -831,8 +831,23 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3u) {
 
     let orig = textureLoad(imageSrc, uid);
     let blur = textureLoad(imageSdw, oid);
-    let shad = settings[0] * blur.a;
+    let shad = settings[1] * blur.a;
     let color = orig + shad * (1.0 - orig.a);
     textureStore(imageTrg, uid.xy, color);
 }
+    
+@compute @workgroup_size(128, 1)
+fn cs_main_fill(@builtin(global_invocation_id) gid: vec3u) {
+    // decode viewport and settings
+    let vmin = vec2u(viewport.xy);
+    let vmax = vec2u(viewport.zw);
+
+    // tex coord
+    let uid = gid.xy + vmin;
+
+    let orig = textureLoad(imageSrc, uid);
+    let fill = settings[0];
+    let color = fill * orig.a * fill.a;
+    textureStore(imageTrg, uid.xy, color);
+}
 )";

src/renderer/wg_engine/tvgWgShaderSrc.h

@@ -46,6 +46,6 @@ extern const char* cShaderSrc_Blit;
 
 // compute shader sources: effects
 extern const char* cShaderSrc_GaussianBlur;
-extern const char* cShaderSrc_DropShadow;
+extern const char* cShaderSrc_Effects;
 
 #endif // _TVG_WG_SHEDER_SRC_H_

src/renderer/wg_engine/tvgWgShaderTypes.cpp

@@ -197,52 +197,53 @@ void WgShaderTypeGradient::updateTexData(const Fill::ColorStop* stops, uint32_t
 }
 
 //************************************************************************
-// WgShaderTypeGaussianBlur
+// WgShaderTypeEffectParams
 //************************************************************************
 
-void WgShaderTypeGaussianBlur::update(float sigma, const Matrix& transform)
+void WgShaderTypeEffectParams::update(const RenderEffectGaussianBlur* gaussian, const Matrix& transform)
-{
-    const float scale = std::sqrt(transform.e11 * transform.e11 + transform.e12 * transform.e12);
-    const float kernel = std::min(WG_GAUSSIAN_KERNEL_SIZE_MAX, 2 * sigma * scale); // kernel size
-    settings[0] = sigma;
-    settings[1] = std::min(WG_GAUSSIAN_KERNEL_SIZE_MAX / kernel, scale);
-    settings[2] = kernel;
-    settings[3] = 0.0f; // unused
-    extend = settings[2] * 2; // kernel
-}
-
-
-void WgShaderTypeGaussianBlur::update(const RenderEffectGaussianBlur* gaussian, const Matrix& transform)
 {
     assert(gaussian);
-    update(gaussian->sigma, transform);
+    const float sigma = gaussian->sigma;
-}
-
-
-void WgShaderTypeGaussianBlur::update(const RenderEffectDropShadow* dropShadow, const Matrix& transform)
-{
-    assert(dropShadow);
-    update(dropShadow->sigma, transform);
-}
-
-//************************************************************************
-// WgShaderTypeDropShadow
-//************************************************************************
-
-void WgShaderTypeDropShadow::update(const RenderEffectDropShadow* dropShadow, const Matrix& transform)
-{
-    assert(dropShadow);
     const float scale = std::sqrt(transform.e11 * transform.e11 + transform.e12 * transform.e12);
+    const float kernel = std::min(WG_GAUSSIAN_KERNEL_SIZE_MAX, 2 * sigma * scale); // kernel size
+    params[0] = sigma;
+    params[1] = std::min(WG_GAUSSIAN_KERNEL_SIZE_MAX / kernel, scale);
+    params[2] = kernel;
+    params[3] = 0.0f;
+    extend = params[2] * 2; // kernel
+}
+
+
+void WgShaderTypeEffectParams::update(const RenderEffectDropShadow* dropShadow, const Matrix& transform)
+{
+    assert(dropShadow);
     const float radian = tvg::deg2rad(90.0f - dropShadow->angle);
+    const float sigma = dropShadow->sigma;
+    const float scale = std::sqrt(transform.e11 * transform.e11 + transform.e12 * transform.e12);
+    const float kernel = std::min(WG_GAUSSIAN_KERNEL_SIZE_MAX, 2 * sigma * scale); // kernel size
     offset = {0, 0};
     if (dropShadow->distance > 0.0f) offset = {
         +1.0f * dropShadow->distance * cosf(radian) * scale,
         -1.0f * dropShadow->distance * sinf(radian) * scale
     };
-    settings[0] = dropShadow->color[0] / 255.0f; // red
+    params[0] = sigma;
-    settings[1] = dropShadow->color[1] / 255.0f; // green
+    params[1] = std::min(WG_GAUSSIAN_KERNEL_SIZE_MAX / kernel, scale);
-    settings[2] = dropShadow->color[2] / 255.0f; // blue
+    params[2] = kernel;
-    settings[3] = dropShadow->color[3] / 255.0f; // alpha
+    params[3] = 0.0f;
-    settings[4] = offset.x;
+    params[4] = dropShadow->color[0] / 255.0f; // red
-    settings[5] = offset.y;
+    params[5] = dropShadow->color[1] / 255.0f; // green
+    params[6] = dropShadow->color[2] / 255.0f; // blue
+    params[7] = dropShadow->color[3] / 255.0f; // alpha
+    params[8] = offset.x;
+    params[9] = offset.y;
+    extend = params[2] * 2; // kernel
+}
+
+
+void WgShaderTypeEffectParams::update(const RenderEffectFill* fill)
+{
+    params[0] = fill->color[0] / 255.0f;
+    params[1] = fill->color[1] / 255.0f;
+    params[2] = fill->color[2] / 255.0f;
+    params[3] = fill->color[3] / 255.0f;
 }

src/renderer/wg_engine/tvgWgShaderTypes.h

@@ -67,25 +67,21 @@ struct WgShaderTypeGradient
     void updateTexData(const Fill::ColorStop* stops, uint32_t stopCnt);
 };
 
-// gaussian settings: sigma, scale, extend
+// gaussian params: sigma, scale, extend
 #define WG_GAUSSIAN_KERNEL_SIZE_MAX (128.0f)
-struct WgShaderTypeGaussianBlur
+// gaussian blur, drop shadow, fill, tint, tritone
+struct WgShaderTypeEffectParams
 {
-    float settings[4]{}; // [0]: sigma, [1]: scale, [2]: kernel size, [3]: unused
+    // gaussian blur: [0]: sigma, [1]: scale, [2]: kernel size
-    uint32_t extend{};
+    // drop shadow:   [0]: sigma, [1]: scale, [2]: kernel size, [4..7]: color, [8, 9]: offset
+    // fill:          [0..3]: color
+    float params[4+4+4]{}; // settings: array<vec4f, 3>;
+    uint32_t extend{};     // gaussian blur extend
+    Point offset{};        // drop shadow offset
 
-    void update(float sigma, const Matrix& transform);
     void update(const RenderEffectGaussianBlur* gaussian, const Matrix& transform);
     void update(const RenderEffectDropShadow* dropShadow, const Matrix& transform);
-};
+    void update(const RenderEffectFill* fill);
-
-// drop shadow settings: color, offset
-struct WgShaderTypeDropShadow
-{
-    float settings[8]{}; // [0..3]: color, [4, 5]: offset
-    Point offset{};
-
-    void update(const RenderEffectDropShadow* dropShadow, const Matrix& transform);
 };
 
 #endif // _TVG_WG_SHADER_TYPES_H_