wg_engine: Introduce drop shadow effect for webgpu renderer

Issue: https://github.com/thorvg/thorvg/issues/3054
Supported color, angle, distance, sigma, quality params

src/renderer/wg_engine/tvgWgCompositor.cpp

@@ -80,8 +80,10 @@ void WgCompositor::releasePools(WgContext& context)
 void WgCompositor::resize(WgContext& context, uint32_t width, uint32_t height) {
     // release existig handles
     if ((this->width != width) || (this->height != height)) {
+        context.layouts.releaseBindGroup(bindGroupStorageTemp);
         // release intermediate render storages
-        storageDstCopy.release(context);
+        storageTemp1.release(context);
+        storageTemp0.release(context);
         // release global stencil buffer handles
         context.releaseTextureView(texViewDepthStencilMS);
         context.releaseTexture(texDepthStencilMS);
@@ -106,7 +108,9 @@ void WgCompositor::resize(WgContext& context, uint32_t width, uint32_t height) {
         texDepthStencilMS = context.createTexAttachement(width, height, WGPUTextureFormat_Depth24PlusStencil8, 4);
         texViewDepthStencilMS = context.createTextureView(texDepthStencilMS);
         // initialize intermediate render storages
-        storageDstCopy.initialize(context, width, height);
+        storageTemp0.initialize(context, width, height);
+        storageTemp1.initialize(context, width, height);
+        bindGroupStorageTemp = context.layouts.createBindGroupStrorage2RO(storageTemp0.texView, storageTemp1.texView);
     }
 }
 
@@ -122,6 +126,25 @@ RenderRegion WgCompositor::shrinkRenderRegion(RenderRegion& rect)
 }
 
 
+void WgCompositor::copyTexture(const WgRenderStorage* dst, const WgRenderStorage* src)
+{
+    const RenderRegion region = { 0, 0, (int32_t)src->width, (int32_t)src->height };
+    copyTexture(dst, src, region);
+}
+
+
+void WgCompositor::copyTexture(const WgRenderStorage* dst, const WgRenderStorage* src, const RenderRegion& region)
+{
+    assert(dst);
+    assert(src);
+    assert(commandEncoder);
+    const WGPUImageCopyTexture texSrc { .texture = src->texture, .origin = { .x = (uint32_t)region.x, .y = (uint32_t)region.y } };
+    const WGPUImageCopyTexture texDst { .texture = dst->texture, .origin = { .x = (uint32_t)region.x, .y = (uint32_t)region.y } };
+    const WGPUExtent3D copySize { .width = (uint32_t)region.w, .height = (uint32_t)region.h, .depthOrArrayLayers = 1 };
+    wgpuCommandEncoderCopyTextureToTexture(commandEncoder, &texSrc, &texDst, &copySize);
+}
+
+
 void WgCompositor::beginRenderPass(WGPUCommandEncoder commandEncoder, WgRenderStorage* target, bool clear, WGPUColor clearColor)
 {
     assert(commandEncoder);
@@ -313,10 +336,7 @@ void WgCompositor::blendShape(WgContext& context, WgRenderDataShape* renderData,
     // copy current render target data to dst storage
     WgRenderStorage *target = currentTarget;
     endRenderPass();
-    const WGPUImageCopyTexture texSrc { .texture = target->texture };
+    copyTexture(&storageTemp0, target);
-    const WGPUImageCopyTexture texDst { .texture = storageDstCopy.texture };
-    const WGPUExtent3D copySize { .width = width, .height = height, .depthOrArrayLayers = 1 };
-    wgpuCommandEncoderCopyTextureToTexture(commandEncoder, &texSrc, &texDst, &copySize);
     beginRenderPass(commandEncoder, target, false);
     // render shape with blend settings
     wgpuRenderPassEncoderSetScissorRect(renderPassEncoder, renderData->viewport.x, renderData->viewport.y, renderData->viewport.w, renderData->viewport.h);
@@ -333,7 +353,7 @@ void WgCompositor::blendShape(WgContext& context, WgRenderDataShape* renderData,
     wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 0);
     wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 0, bindGroupViewMat, 0, nullptr);
     wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 1, renderData->bindGroupPaint, 0, nullptr);
-    wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 3, storageDstCopy.bindGroupTexure, 0, nullptr);
+    wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 3, storageTemp0.bindGroupTexure, 0, nullptr);
     uint32_t blendMethodInd = (uint32_t)blendMethod;
     WgRenderSettings& settings = renderData->renderSettingsShape;
     if (settings.fillType == WgRenderSettingsType::Solid) {
@@ -444,10 +464,7 @@ void WgCompositor::blendStrokes(WgContext& context, WgRenderDataShape* renderDat
     // copy current render target data to dst storage
     WgRenderStorage *target = currentTarget;
     endRenderPass();
-    const WGPUImageCopyTexture texSrc { .texture = target->texture };
+    copyTexture(&storageTemp0, target);
-    const WGPUImageCopyTexture texDst { .texture = storageDstCopy.texture };
-    const WGPUExtent3D copySize { .width = width, .height = height, .depthOrArrayLayers = 1 };
-    wgpuCommandEncoderCopyTextureToTexture(commandEncoder, &texSrc, &texDst, &copySize);
     beginRenderPass(commandEncoder, target, false);
     wgpuRenderPassEncoderSetScissorRect(renderPassEncoder, renderData->viewport.x, renderData->viewport.y, renderData->viewport.w, renderData->viewport.h);
     // draw strokes to stencil (first pass)
@@ -463,7 +480,7 @@ void WgCompositor::blendStrokes(WgContext& context, WgRenderDataShape* renderDat
         wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 0);
         wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 0, bindGroupViewMat, 0, nullptr);
         wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 1, renderData->bindGroupPaint, 0, nullptr);
-        wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 3, storageDstCopy.bindGroupTexure, 0, nullptr);
+        wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 3, storageTemp0.bindGroupTexure, 0, nullptr);
         uint32_t blendMethodInd = (uint32_t)blendMethod;
         WgRenderSettings& settings = renderData->renderSettingsStroke;
         if (settings.fillType == WgRenderSettingsType::Solid) {
@@ -557,10 +574,7 @@ void WgCompositor::blendImage(WgContext& context, WgRenderDataPicture* renderDat
     // copy current render target data to dst storage
     WgRenderStorage *target = currentTarget;
     endRenderPass();
-    const WGPUImageCopyTexture texSrc { .texture = target->texture };
+    copyTexture(&storageTemp0, target);
-    const WGPUImageCopyTexture texDst { .texture = storageDstCopy.texture };
-    const WGPUExtent3D copySize { .width = width, .height = height, .depthOrArrayLayers = 1 };
-    wgpuCommandEncoderCopyTextureToTexture(commandEncoder, &texSrc, &texDst, &copySize);
     beginRenderPass(commandEncoder, target, false);
     // setup stencil rules
     wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 255);
@@ -574,7 +588,7 @@ void WgCompositor::blendImage(WgContext& context, WgRenderDataPicture* renderDat
     wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 0, bindGroupViewMat, 0, nullptr);
     wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 1, renderData->bindGroupPaint, 0, nullptr);
     wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 2, renderData->bindGroupPicture, 0, nullptr);
-    wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 3, storageDstCopy.bindGroupTexure, 0, nullptr);
+    wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 3, storageTemp0.bindGroupTexure, 0, nullptr);
     wgpuRenderPassEncoderSetPipeline(renderPassEncoder, pipelines.image_blend[blendMethodInd]);
     renderData->meshData.drawImage(context, renderPassEncoder);
 };
@@ -631,10 +645,7 @@ void WgCompositor::blendScene(WgContext& context, WgRenderStorage* scene, WgComp
     // copy current render target data to dst storage
     WgRenderStorage *target = currentTarget;
     endRenderPass();
-    const WGPUImageCopyTexture texSrc { .texture = target->texture };
+    copyTexture(&storageTemp0, target);
-    const WGPUImageCopyTexture texDst { .texture = storageDstCopy.texture };
-    const WGPUExtent3D copySize { .width = width, .height = height, .depthOrArrayLayers = 1 };
-    wgpuCommandEncoderCopyTextureToTexture(commandEncoder, &texSrc, &texDst, &copySize);
     beginRenderPass(commandEncoder, target, false);
     // blend scene
     uint32_t blendMethodInd = (uint32_t)compose->blend;
@@ -642,7 +653,7 @@ void WgCompositor::blendScene(WgContext& context, WgRenderStorage* scene, WgComp
     wgpuRenderPassEncoderSetScissorRect(renderPassEncoder, rect.x, rect.y, rect.w, rect.h);
     wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 0);
     wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 0, scene->bindGroupTexure, 0, nullptr);
-    wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 1, storageDstCopy.bindGroupTexure, 0, nullptr);
+    wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 1, storageTemp0.bindGroupTexure, 0, nullptr);
     wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 2, bindGroupOpacities[compose->opacity], 0, nullptr);
     wgpuRenderPassEncoderSetPipeline(renderPassEncoder, pipelines.scene_blend[blendMethodInd]);
     meshData.drawImage(context, renderPassEncoder);
@@ -739,7 +750,7 @@ void WgCompositor::clearClipPath(WgContext& context, WgRenderDataPaint* paint)
 }
 
 
-void WgCompositor::gaussianBlur(WgContext& context, WgRenderStorage* dst, const RenderEffectGaussianBlur* params, const WgCompose* compose)
+bool WgCompositor::gaussianBlur(WgContext& context, WgRenderStorage* dst, const RenderEffectGaussianBlur* params, const WgCompose* compose)
 {
     assert(dst);
     assert(params);
@@ -750,7 +761,7 @@ void WgCompositor::gaussianBlur(WgContext& context, WgRenderStorage* dst, const
     auto aabb = compose->aabb;
     auto viewport = compose->rdViewport;
     WgRenderStorage* sbuff = dst;
-    WgRenderStorage* dbuff = &storageDstCopy;
+    WgRenderStorage* dbuff = &storageTemp0;
 
     // begin compute pass
     WGPUComputePassDescriptor computePassDesc{ .label = "Compute pass gaussian blur" };
@@ -782,10 +793,64 @@ void WgCompositor::gaussianBlur(WgContext& context, WgRenderStorage* dst, const
     wgpuComputePassEncoderRelease(computePassEncoder);
 
     // if final result stored in intermidiate buffer we must copy result to destination buffer
-    if (sbuff == &storageDstCopy) {
+    if (sbuff == &storageTemp0) 
-        const WGPUImageCopyTexture texSrc { .texture = sbuff->texture, .origin = { .x = (uint32_t)aabb.x, .y = (uint32_t)aabb.y } };
+        copyTexture(sbuff, dbuff, aabb);
-        const WGPUImageCopyTexture texDst { .texture = dbuff->texture, .origin = { .x = (uint32_t)aabb.x, .y = (uint32_t)aabb.y } };
+
-        const WGPUExtent3D copySize { .width = (uint32_t)aabb.w, .height = (uint32_t)aabb.h, .depthOrArrayLayers = 1 };
+    return true;
-        wgpuCommandEncoderCopyTextureToTexture(commandEncoder, &texSrc, &texDst, &copySize);
+}
+
+
+bool WgCompositor::dropShadow(WgContext& context, WgRenderStorage* dst, const RenderEffectDropShadow* params, const WgCompose* compose)
+{
+    assert(dst);
+    assert(params);
+    assert(params->rd);
+    assert(compose->rdViewport);
+    assert(!renderPassEncoder);
+
+    auto renderDataDropShadow = (WgRenderDataDropShadow*)params->rd;
+    auto aabb = compose->aabb;
+    auto viewport = compose->rdViewport;
+
+    { // apply blur
+        copyTexture(&storageTemp1, dst, aabb);
+        WgRenderStorage* sbuff = &storageTemp1;
+        WgRenderStorage* dbuff = &storageTemp0;
+        WGPUComputePassDescriptor computePassDesc{ .label = "Compute pass drop shadow blur" };
+        WGPUComputePassEncoder computePassEncoder = wgpuCommandEncoderBeginComputePass(commandEncoder, &computePassDesc);
+        // 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, 3, viewport->bindGroupViewport, 0, nullptr);
+        wgpuComputePassEncoderSetPipeline(computePassEncoder, pipelines.gaussian_horz);
+        wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, (aabb.w - 1) / 128 + 1, aabb.h, 1);
+        std::swap(sbuff, dbuff);
+        // 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, 3, viewport->bindGroupViewport, 0, nullptr);
+        wgpuComputePassEncoderSetPipeline(computePassEncoder, pipelines.gaussian_vert);
+        wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, aabb.w, (aabb.h - 1) / 128 + 1, 1);
+        std::swap(sbuff, dbuff);
+        wgpuComputePassEncoderEnd(computePassEncoder);
+        wgpuComputePassEncoderRelease(computePassEncoder);
     }
+    
+    { // blend origin (temp0), shadow (temp1) to destination
+        copyTexture(&storageTemp0, dst, aabb);
+        WGPUComputePassDescriptor computePassDesc{ .label = "Compute pass drop shadow blend" };
+        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, 3, viewport->bindGroupViewport, 0, nullptr);
+        wgpuComputePassEncoderSetPipeline(computePassEncoder, pipelines.dropshadow);
+        wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, (aabb.w - 1) / 128 + 1, aabb.h, 1);
+        wgpuComputePassEncoderEnd(computePassEncoder);
+        wgpuComputePassEncoderRelease(computePassEncoder);
+    }
+
+    return true;
 }

src/renderer/wg_engine/tvgWgCompositor.h

@@ -54,7 +54,9 @@ private:
     WGPUCommandEncoder commandEncoder{};
     WgRenderStorage* currentTarget{};
     // intermediate render storages
-    WgRenderStorage storageDstCopy;
+    WgRenderStorage storageTemp0;
+    WgRenderStorage storageTemp1;
+    WGPUBindGroup bindGroupStorageTemp{};
     // composition and blend geometries
     WgMeshData meshData;
     // render target dimensions
@@ -63,6 +65,9 @@ private:
     
     // viewport utilities
     RenderRegion shrinkRenderRegion(RenderRegion& rect);
+    void copyTexture(const WgRenderStorage* dst, const WgRenderStorage* src);
+    void copyTexture(const WgRenderStorage* dst, const WgRenderStorage* src, const RenderRegion& region);
+
 
     // shapes
     void drawShape(WgContext& context, WgRenderDataShape* renderData);
@@ -107,7 +112,8 @@ public:
     void blit(WgContext& context, WGPUCommandEncoder encoder, WgRenderStorage* src, WGPUTextureView dstView);
 
     // effects
-    void gaussianBlur(WgContext& context, WgRenderStorage* dst, const RenderEffectGaussianBlur* params, const WgCompose* compose);
+    bool gaussianBlur(WgContext& context, WgRenderStorage* dst, const RenderEffectGaussianBlur* params, const WgCompose* compose);
+    bool dropShadow(WgContext& context, WgRenderStorage* dst, const RenderEffectDropShadow* params, const WgCompose* compose);
 };
 
 #endif // _TVG_WG_COMPOSITOR_H_

src/renderer/wg_engine/tvgWgPipelines.cpp

@@ -185,6 +185,7 @@ void WgPipelines::initialize(WgContext& context)
     const WGPUBindGroupLayout bindGroupLayoutsBlit[] { layouts.layoutTexSampled };
     // bind group layouts effects
     const WGPUBindGroupLayout bindGroupLayoutsGaussian[] { layouts.layoutTexStrorage1RO, layouts.layoutTexStrorage1WO, layouts.layoutBuffer1Un, layouts.layoutBuffer1Un };
+    const WGPUBindGroupLayout bindGroupLayoutsDropShadow[] { 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,6 +223,7 @@ void WgPipelines::initialize(WgContext& context)
     shader_blit = createShaderModule(context.device, "The shader blit", cShaderSrc_Blit);
     // shader effects
     shader_gaussian = createShaderModule(context.device, "The shader gaussian", cShaderSrc_GaussianBlur);
+    shader_dropshadow = createShaderModule(context.device, "The shader drop shadow", cShaderSrc_DropShadow);
 
     // layouts
     layout_stencil = createPipelineLayout(context.device, bindGroupLayoutsStencil, 2);
@@ -242,6 +244,7 @@ void WgPipelines::initialize(WgContext& context)
     layout_blit = createPipelineLayout(context.device, bindGroupLayoutsBlit, 1);
     // layout effects
     layout_gaussian = createPipelineLayout(context.device, bindGroupLayoutsGaussian, 4);
+    layout_dropshadow = createPipelineLayout(context.device, bindGroupLayoutsDropShadow, 4);
 
     // render pipeline nonzero
     nonzero = createRenderPipeline(
@@ -447,17 +450,16 @@ void WgPipelines::initialize(WgContext& context)
         depthStencilStateScene, multisampleStateX1);
 
     // compute pipeline gaussian blur
-    gaussian_horz = createComputePipeline(
+    gaussian_horz = createComputePipeline(context.device, "The compute pipeline gaussian blur horizontal", shader_gaussian, "cs_main_horz", layout_gaussian);
-        context.device, "The compute pipeline gaussian blur horizontal",
+    gaussian_vert = createComputePipeline(context.device, "The compute pipeline gaussian blur vertical",   shader_gaussian, "cs_main_vert", layout_gaussian);
-        shader_gaussian, "cs_main_horz", layout_gaussian);
+    // compute pipeline drop shadow
-    gaussian_vert = createComputePipeline(
+    dropshadow = createComputePipeline(context.device, "The compute pipeline drop shadow blend", shader_dropshadow, "cs_main", layout_dropshadow);
-        context.device, "The compute pipeline gaussian blur vertical",
-        shader_gaussian, "cs_main_vert", layout_gaussian);
 }
 
 void WgPipelines::releaseGraphicHandles(WgContext& context)
 {
     // pipeline effects
+    releaseComputePipeline(dropshadow);
     releaseComputePipeline(gaussian_vert);
     releaseComputePipeline(gaussian_horz);
     // pipeline blit
@@ -490,6 +492,7 @@ void WgPipelines::releaseGraphicHandles(WgContext& context)
     releaseRenderPipeline(evenodd);
     releaseRenderPipeline(nonzero);
     // layouts
+    releasePipelineLayout(layout_dropshadow);
     releasePipelineLayout(layout_gaussian);
     releasePipelineLayout(layout_blit);
     releasePipelineLayout(layout_scene_compose);
@@ -504,6 +507,7 @@ void WgPipelines::releaseGraphicHandles(WgContext& context)
     releasePipelineLayout(layout_depth);
     releasePipelineLayout(layout_stencil);
     // shaders
+    releaseShaderModule(shader_dropshadow);
     releaseShaderModule(shader_gaussian);
     releaseShaderModule(shader_blit);
     releaseShaderModule(shader_scene_compose);

src/renderer/wg_engine/tvgWgPipelines.h

@@ -48,6 +48,7 @@ private:
     WGPUShaderModule shader_blit{};
     // shader effects
     WGPUShaderModule shader_gaussian;
+    WGPUShaderModule shader_dropshadow;
 
     // layouts helpers
     WGPUPipelineLayout layout_stencil{};
@@ -68,6 +69,7 @@ private:
     WGPUPipelineLayout layout_blit{};
     // layouts effects
     WGPUPipelineLayout layout_gaussian{};
+    WGPUPipelineLayout layout_dropshadow{};
 public:
     // pipelines stencil markup
     WGPURenderPipeline nonzero{};
@@ -98,6 +100,7 @@ public:
     // effects
     WGPUComputePipeline gaussian_horz{};
     WGPUComputePipeline gaussian_vert{};
+    WGPUComputePipeline dropshadow{};
 private:
     void releaseGraphicHandles(WgContext& context);
     WGPUShaderModule createShaderModule(WGPUDevice device, const char* label, const char* code);

src/renderer/wg_engine/tvgWgRenderData.cpp

@@ -640,3 +640,78 @@ void WgRenderDataGaussianPool::release(WgContext& context)
     mPool.clear();
     mList.clear();
 }
+
+//***********************************************************************
+// WgRenderDataDropShadow
+//***********************************************************************
+
+void WgRenderDataDropShadow::update(WgContext& context, RenderEffectDropShadow* dropShadow, const Matrix& transform)
+{
+    assert(dropShadow);
+    // compute gaussian blur data
+    WgShaderTypeGaussianBlur gaussianSettings;
+    gaussianSettings.update(dropShadow, transform);
+    // 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;
+    offset = dropShadowSettings.offset;
+}
+
+
+void WgRenderDataDropShadow::release(WgContext& context)
+{
+    context.releaseBuffer(bufferSettings);
+    context.releaseBuffer(bufferGaussian);
+    context.layouts.releaseBindGroup(bindGroupDropShadow);
+    context.layouts.releaseBindGroup(bindGroupGaussian);
+}
+
+//***********************************************************************
+// WgRenderDataGaussianPool
+//***********************************************************************
+
+WgRenderDataDropShadow* WgRenderDataDropShadowPool::allocate(WgContext& context)
+{
+    WgRenderDataDropShadow* renderData{};
+    if (mPool.count > 0) {
+        renderData = mPool.last();
+        mPool.pop();
+    } else {
+        renderData = new WgRenderDataDropShadow();
+        mList.push(renderData);
+    }
+    return renderData;
+}
+
+
+void WgRenderDataDropShadowPool::free(WgContext& context, WgRenderDataDropShadow* renderData)
+{
+    if (renderData) mPool.push(renderData);
+}
+
+
+void WgRenderDataDropShadowPool::release(WgContext& context)
+{
+    ARRAY_FOREACH(p, mList) {
+        (*p)->release(context);
+        delete(*p);
+    }
+    mPool.clear();
+    mList.clear();
+}

src/renderer/wg_engine/tvgWgRenderData.h

@@ -215,4 +215,31 @@ public:
     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 release(WgContext& context);
+};
+
+class WgRenderDataDropShadowPool {
+private:
+    // pool contains all created but unused render data for drop shadow
+    Array<WgRenderDataDropShadow*> mPool;
+    // list contains all created render data for drop shadow
+    // to ensure that all created instances will be released
+    Array<WgRenderDataDropShadow*> mList;
+public:
+    WgRenderDataDropShadow* allocate(WgContext& context);
+    void free(WgContext& context, WgRenderDataDropShadow* renderData);
+    void release(WgContext& context);
+};
+
 #endif // _TVG_WG_RENDER_DATA_H_

src/renderer/wg_engine/tvgWgRenderer.cpp

@@ -53,6 +53,7 @@ void WgRenderer::release()
     mRenderDataShapePool.release(mContext);
     mRenderDataPicturePool.release(mContext);
     mRenderDataGaussianPool.release(mContext);
+    mRenderDataDropShadowPool.release(mContext);
     mRenderDataViewportPool.release(mContext);
     WgMeshDataPool::gMeshDataPool->release(mContext);
 
@@ -531,12 +532,24 @@ void WgRenderer::prepare(RenderEffect* effect, const Matrix& transform)
         }
         renderDataGaussian->update(mContext, gaussianBlur, transform);
         effect->valid = true;
+    } else
+    // prepare drop shadow data
+    if (effect->type == SceneEffect::DropShadow) {
+        auto dropShadow = (RenderEffectDropShadow*)effect;
+        auto renderDataDropShadow = (WgRenderDataDropShadow*)dropShadow->rd;
+        if (!renderDataDropShadow) {
+            renderDataDropShadow = mRenderDataDropShadowPool.allocate(mContext);
+            dropShadow->rd = renderDataDropShadow;
+        }
+        renderDataDropShadow->update(mContext, dropShadow, transform);
+        effect->valid = true;
     }
 }
 
 
 bool WgRenderer::region(RenderEffect* effect)
 {
+    // update gaussian blur region
     if (effect->type == SceneEffect::GaussianBlur) {
         auto gaussian = (RenderEffectGaussianBlur*)effect;
         auto renderDataGaussian = (WgRenderDataGaussian*)gaussian->rd;
@@ -549,6 +562,16 @@ bool WgRenderer::region(RenderEffect* effect)
             gaussian->extend.h = +renderDataGaussian->extend * 2;
         }
         return true;
+    } else
+    // update drop shadow region
+    if (effect->type == SceneEffect::DropShadow) {
+        auto dropShadow = (RenderEffectDropShadow*)effect;
+        auto renderDataDropShadow = (WgRenderDataDropShadow*)dropShadow->rd;
+        dropShadow->extend.x = -(renderDataDropShadow->extend + std::abs(renderDataDropShadow->offset.x));
+        dropShadow->extend.w = +(renderDataDropShadow->extend + std::abs(renderDataDropShadow->offset.x)) * 2;
+        dropShadow->extend.y = -(renderDataDropShadow->extend + std::abs(renderDataDropShadow->offset.y));
+        dropShadow->extend.h = +(renderDataDropShadow->extend + std::abs(renderDataDropShadow->offset.y)) * 2;
+        return true;
     }
     return false;
 }
@@ -558,14 +581,13 @@ bool WgRenderer::render(RenderCompositor* cmp, const RenderEffect* effect, TVG_U
 {
     // we must to end current render pass to resolve ms texture before effect
     mCompositor.endRenderPass();
+    WgCompose* comp = (WgCompose*)cmp;
+    WgRenderStorage* dst = mRenderStorageStack.last();
 
-    // handle gaussian blur
+    switch (effect->type) {
-    if (effect->type == SceneEffect::GaussianBlur) {
+        case SceneEffect::GaussianBlur: return mCompositor.gaussianBlur(mContext, dst, (RenderEffectGaussianBlur*)effect, comp);
-        WgCompose* comp = (WgCompose*)cmp;
+        case SceneEffect::DropShadow: return mCompositor.dropShadow(mContext, dst, (RenderEffectDropShadow*)effect, comp);
-        WgRenderStorage* dst = mRenderStorageStack.last();
+        default: return false;
-        RenderEffectGaussianBlur* gaussianBlur = (RenderEffectGaussianBlur*)effect;
-        mCompositor.gaussianBlur(mContext, dst, gaussianBlur, comp);
-        return true;
     }
     return false;
 }
@@ -573,12 +595,12 @@ bool WgRenderer::render(RenderCompositor* cmp, const RenderEffect* effect, TVG_U
 
 void WgRenderer::dispose(RenderEffect* effect)
 {
-    // dispose gaussian blur data
+    switch (effect->type) {
-    if (effect->type == SceneEffect::GaussianBlur) {
+        case SceneEffect::GaussianBlur: mRenderDataGaussianPool.free(mContext, (WgRenderDataGaussian*)effect->rd);
-        auto gaussianBlur = (RenderEffectGaussianBlur*)effect;
+        case SceneEffect::DropShadow: mRenderDataDropShadowPool.free(mContext, (WgRenderDataDropShadow*)effect->rd);
-        mRenderDataGaussianPool.free(mContext, (WgRenderDataGaussian*)gaussianBlur->rd);
+        default: effect->rd = nullptr;
-        gaussianBlur->rd = nullptr;
     }
+    effect->rd = nullptr;
 };
 
 

src/renderer/wg_engine/tvgWgRenderer.h

@@ -86,6 +86,7 @@ private:
     WgRenderDataShapePool mRenderDataShapePool;
     WgRenderDataPicturePool mRenderDataPicturePool;
     WgRenderDataGaussianPool mRenderDataGaussianPool;
+    WgRenderDataDropShadowPool mRenderDataDropShadowPool;
     WgRenderDataViewportPool mRenderDataViewportPool;
 
     // rendering context

src/renderer/wg_engine/tvgWgShaderSrc.cpp

@@ -708,19 +708,6 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f {
 // compute shader source: merge clip path masks
 //************************************************************************
 
-const char* cShaderSrc_MergeMasks = R"(
-@group(0) @binding(0) var imageMsk0 : texture_storage_2d<rgba8unorm, read>;
-@group(1) @binding(0) var imageMsk1 : texture_storage_2d<rgba8unorm, read>;
-@group(2) @binding(0) var imageTrg  : texture_storage_2d<rgba8unorm, write>;
-
-@compute @workgroup_size(8, 8)
-fn cs_main(@builtin(global_invocation_id) id: vec3u) {
-    let colorMsk0 = textureLoad(imageMsk0, id.xy);
-    let colorMsk1 = textureLoad(imageMsk1, id.xy);
-    textureStore(imageTrg, id.xy, colorMsk0 * colorMsk1);
-}
-)";
-
 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>;
@@ -821,6 +808,31 @@ fn cs_main_vert(@builtin(global_invocation_id) gid: vec3u,
 
     // store result
     textureStore(imageDst, uid, color / sum);
-    //textureStore(imageDst, uid, vec4f(1.0, 0.0, 0.0, 1.0));
+}
+)";
+
+const char* cShaderSrc_DropShadow = 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(3) @binding(0) var<uniform> viewport: vec4f;
+
+@compute @workgroup_size(128, 1)
+fn cs_main(@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);
+
+    // tex coord
+    let uid = gid.xy + vmin;
+    let oid = clamp(vec2u(vec2i(uid) - voff), vmin, vmax);
+
+    let orig = textureLoad(imageSrc, uid);
+    let blur = textureLoad(imageSdw, oid);
+    let shad = settings[0] * blur.a;
+    let color = orig + shad * (1.0 - orig.a);
+    textureStore(imageTrg, uid.xy, color);
 }
 )";

src/renderer/wg_engine/tvgWgShaderSrc.h

@@ -45,7 +45,7 @@ extern const char* cShaderSrc_Scene_Compose;
 extern const char* cShaderSrc_Blit;
 
 // compute shader sources: effects
-extern const char* cShaderSrc_MergeMasks;
 extern const char* cShaderSrc_GaussianBlur;
+extern const char* cShaderSrc_DropShadow;
 
 #endif // _TVG_WG_SHEDER_SRC_H_

src/renderer/wg_engine/tvgWgShaderTypes.cpp

@@ -200,14 +200,49 @@ void WgShaderTypeGradient::updateTexData(const Fill::ColorStop* stops, uint32_t
 // WgShaderTypeGaussianBlur
 //************************************************************************
 
-void WgShaderTypeGaussianBlur::update(const RenderEffectGaussianBlur* gaussian, const Matrix& transform)
+void WgShaderTypeGaussianBlur::update(float sigma, const Matrix& transform)
 {
-    assert(gaussian);
-    const float sigma = gaussian->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
     settings[0] = sigma;
     settings[1] = std::min(WG_GAUSSIAN_KERNEL_SIZE_MAX / kernel, scale);
     settings[2] = kernel;
-    extend = settings[2] * 2;
+    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);
+}
+
+
+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 radian = tvg::deg2rad(90.0f - dropShadow->angle);
+    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
+    settings[1] = dropShadow->color[1] / 255.0f; // green
+    settings[2] = dropShadow->color[2] / 255.0f; // blue
+    settings[3] = dropShadow->color[3] / 255.0f; // alpha
+    settings[4] = offset.x;
+    settings[5] = offset.y;
 }

src/renderer/wg_engine/tvgWgShaderTypes.h

@@ -74,7 +74,18 @@ struct WgShaderTypeGaussianBlur
     float settings[4]{}; // [0]: sigma, [1]: scale, [2]: kernel size, [3]: unused
     uint32_t extend{};
 
+    void update(float sigma, const Matrix& transform);
     void update(const RenderEffectGaussianBlur* gaussian, const Matrix& transform);
+    void update(const RenderEffectDropShadow* dropShadow, const Matrix& transform);
+};
+
+// 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_