wg_engine: ClipPath support

[issues 1479: ClipPath](#1479)

Supports ClipPath composition.
Clip path composition is an only composition type who doesn't ignore blend method.
Clip path is a combination of composition approach and blend approach using compute shader

src/renderer/wg_engine/tvgWgBindGroups.cpp

@@ -36,11 +36,15 @@ WGPUBindGroupLayout WgBindGroupTexture::layout = nullptr;
 WGPUBindGroupLayout WgBindGroupTextureStorageRgba::layout = nullptr;
 WGPUBindGroupLayout WgBindGroupTextureStorageBgra::layout = nullptr;
 WGPUBindGroupLayout WgBindGroupTextureSampled::layout = nullptr;
+WGPUBindGroupLayout WgBindGroupTexMaskCompose::layout = nullptr;
 WGPUBindGroupLayout WgBindGroupTexComposeBlend::layout = nullptr;
 WGPUBindGroupLayout WgBindGroupOpacity::layout = nullptr;
 WGPUBindGroupLayout WgBindGroupBlendMethod::layout = nullptr;
 WGPUBindGroupLayout WgBindGroupCompositeMethod::layout = nullptr;
 
+///////////////////////////////////////////////////////////////////////////////
+// WgBindGroupCanvas
+///////////////////////////////////////////////////////////////////////////////
 
 WGPUBindGroupLayout WgBindGroupCanvas::getLayout(WGPUDevice device)
 {
@@ -78,6 +82,9 @@ void WgBindGroupCanvas::release()
     releaseBuffer(uBufferViewMat);
 }
 
+///////////////////////////////////////////////////////////////////////////////
+// WgBindGroupPaint
+///////////////////////////////////////////////////////////////////////////////
 
 WGPUBindGroupLayout WgBindGroupPaint::getLayout(WGPUDevice device)
 {
@@ -123,6 +130,9 @@ void WgBindGroupPaint::release()
     releaseBuffer(uBufferModelMat);
 }
 
+///////////////////////////////////////////////////////////////////////////////
+// WgBindGroupSolidColor
+///////////////////////////////////////////////////////////////////////////////
 
 WGPUBindGroupLayout WgBindGroupSolidColor::getLayout(WGPUDevice device)
 {
@@ -163,6 +173,9 @@ void WgBindGroupSolidColor::release()
     releaseBuffer(uBufferSolidColor);
 }
 
+///////////////////////////////////////////////////////////////////////////////
+// WgBindGroupLinearGradient
+///////////////////////////////////////////////////////////////////////////////
 
 WGPUBindGroupLayout WgBindGroupLinearGradient::getLayout(WGPUDevice device)
 {
@@ -203,6 +216,9 @@ void WgBindGroupLinearGradient::release()
     releaseBuffer(uBufferLinearGradient);
 }
 
+///////////////////////////////////////////////////////////////////////////////
+// WgBindGroupRadialGradient
+///////////////////////////////////////////////////////////////////////////////
 
 WGPUBindGroupLayout WgBindGroupRadialGradient::getLayout(WGPUDevice device)
 {
@@ -243,6 +259,9 @@ void WgBindGroupRadialGradient::release()
     releaseBindGroup(mBindGroup);
 }
 
+///////////////////////////////////////////////////////////////////////////////
+// WgBindGroupPicture
+///////////////////////////////////////////////////////////////////////////////
 
 WGPUBindGroupLayout WgBindGroupPicture::getLayout(WGPUDevice device)
 {
@@ -280,6 +299,9 @@ void WgBindGroupPicture::release()
     releaseBindGroup(mBindGroup);
 }
 
+///////////////////////////////////////////////////////////////////////////////
+// WgBindGroupTexture
+///////////////////////////////////////////////////////////////////////////////
 
 WGPUBindGroupLayout WgBindGroupTexture::getLayout(WGPUDevice device)
 {
@@ -315,6 +337,9 @@ void WgBindGroupTexture::release()
     releaseBindGroup(mBindGroup);
 }
 
+///////////////////////////////////////////////////////////////////////////////
+// WgBindGroupTextureStorageRgba
+///////////////////////////////////////////////////////////////////////////////
 
 WGPUBindGroupLayout WgBindGroupTextureStorageRgba::getLayout(WGPUDevice device)
 {
@@ -350,6 +375,9 @@ void WgBindGroupTextureStorageRgba::release()
     releaseBindGroup(mBindGroup);
 }
 
+///////////////////////////////////////////////////////////////////////////////
+// WgBindGroupTextureStorageBgra
+///////////////////////////////////////////////////////////////////////////////
 
 WGPUBindGroupLayout WgBindGroupTextureStorageBgra::getLayout(WGPUDevice device)
 {
@@ -385,6 +413,9 @@ void WgBindGroupTextureStorageBgra::release()
     releaseBindGroup(mBindGroup);
 }
 
+///////////////////////////////////////////////////////////////////////////////
+// WgBindGroupTextureSampled
+///////////////////////////////////////////////////////////////////////////////
 
 WGPUBindGroupLayout WgBindGroupTextureSampled::getLayout(WGPUDevice device)
 {
@@ -422,6 +453,9 @@ void WgBindGroupTextureSampled::release()
     releaseBindGroup(mBindGroup);
 }
 
+///////////////////////////////////////////////////////////////////////////////
+// WgBindGroupTexComposeBlend
+///////////////////////////////////////////////////////////////////////////////
 
 WGPUBindGroupLayout WgBindGroupTexComposeBlend::getLayout(WGPUDevice device)
 {
@@ -461,7 +495,49 @@ void WgBindGroupTexComposeBlend::release()
     releaseBindGroup(mBindGroup);
 }
 
+///////////////////////////////////////////////////////////////////////////////
+// WgBindGroupTexMaskCompose
+///////////////////////////////////////////////////////////////////////////////
 
+WGPUBindGroupLayout WgBindGroupTexMaskCompose::getLayout(WGPUDevice device)
+{
+    if (layout) return layout;
+    const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] {
+        makeBindGroupLayoutEntryStorage(0, WGPUStorageTextureAccess_ReadOnly, WGPUTextureFormat_RGBA8Unorm),
+        makeBindGroupLayoutEntryStorage(1, WGPUStorageTextureAccess_ReadWrite, WGPUTextureFormat_RGBA8Unorm)
+    };
+    layout = createBindGroupLayout(device, bindGroupLayoutEntries, 2);
+    assert(layout);
+    return layout;
+}
+
+
+void WgBindGroupTexMaskCompose::releaseLayout()
+{
+    releaseBindGroupLayout(layout);
+}
+
+
+void WgBindGroupTexMaskCompose::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexMsk0, WGPUTextureView uTexMsk1)
+{
+    release();
+    const WGPUBindGroupEntry bindGroupEntries[] {
+        makeBindGroupEntryTextureView(0, uTexMsk0),
+        makeBindGroupEntryTextureView(1, uTexMsk1),
+    };
+    mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 2);
+    assert(mBindGroup);
+}
+
+
+void WgBindGroupTexMaskCompose::release()
+{
+    releaseBindGroup(mBindGroup);
+}
+
+///////////////////////////////////////////////////////////////////////////////
+// WgBindGroupOpacity
+///////////////////////////////////////////////////////////////////////////////
 
 WGPUBindGroupLayout WgBindGroupOpacity::getLayout(WGPUDevice device)
 {
@@ -500,6 +576,9 @@ void WgBindGroupOpacity::release()
     releaseBindGroup(mBindGroup);
 }
 
+///////////////////////////////////////////////////////////////////////////////
+// WgBindGroupBlendMethod
+///////////////////////////////////////////////////////////////////////////////
 
 WGPUBindGroupLayout WgBindGroupBlendMethod::getLayout(WGPUDevice device)
 {
@@ -538,6 +617,9 @@ void WgBindGroupBlendMethod::release()
     releaseBindGroup(mBindGroup);
 }
 
+///////////////////////////////////////////////////////////////////////////////
+// WgBindGroupCompositeMethod
+///////////////////////////////////////////////////////////////////////////////
 
 WGPUBindGroupLayout WgBindGroupCompositeMethod::getLayout(WGPUDevice device)
 {

src/renderer/wg_engine/tvgWgBindGroups.h

@@ -170,6 +170,20 @@ struct WgBindGroupTexComposeBlend : public WgBindGroup
 };
 
 
+// @group(0)
+struct WgBindGroupTexMaskCompose : public WgBindGroup
+{
+    static WGPUBindGroupLayout layout;
+    static WGPUBindGroupLayout getLayout(WGPUDevice device);
+    static void releaseLayout();
+
+    void initialize(WGPUDevice device, WGPUQueue queue,
+                    WGPUTextureView uTexMsk0,
+                    WGPUTextureView uTexMsk1);
+    void release();
+};
+
+
 // @group(1 or 2)
 struct WgBindGroupOpacity : public WgBindGroup
 {

src/renderer/wg_engine/tvgWgPipelines.cpp

@@ -49,7 +49,7 @@ void WgPipelineFillShapeWinding::initialize(WGPUDevice device)
     WGPUCompareFunction stencilFunctionBack = WGPUCompareFunction_Always;
     WGPUStencilOperation stencilOperationBack = WGPUStencilOperation_DecrementWrap;
 
-    // sheder source and labels
+    // shader source and labels
     auto shaderSource = cShaderSource_PipelineFill;
     auto shaderLabel = "The shader fill";
     auto pipelineLabel = "The render pipeline fill shape winding";
@@ -83,7 +83,7 @@ void WgPipelineFillShapeEvenOdd::initialize(WGPUDevice device)
     WGPUCompareFunction stencilFunctionBack = WGPUCompareFunction_Always;
     WGPUStencilOperation stencilOperationBack = WGPUStencilOperation_Invert;
 
-    // sheder source and labels
+    // shader source and labels
     auto shaderSource = cShaderSource_PipelineFill;
     auto shaderLabel = "The shader fill";
     auto pipelineLabel = "The render pipeline fill shape Even Odd";
@@ -115,7 +115,7 @@ void WgPipelineFillStroke::initialize(WGPUDevice device)
     WGPUCompareFunction stencilFunction = WGPUCompareFunction_Always;
     WGPUStencilOperation stencilOperation = WGPUStencilOperation_Replace;
 
-    // sheder source and labels
+    // shader source and labels
     auto shaderSource = cShaderSource_PipelineFill;
     auto shaderLabel = "The shader fill";
     auto pipelineLabel = "The render pipeline fill stroke";
@@ -129,6 +129,38 @@ void WgPipelineFillStroke::initialize(WGPUDevice device)
 }
 
 
+void WgPipelineClipMask::initialize(WGPUDevice device)
+{
+    // vertex and buffers settings
+    WGPUVertexAttribute vertexAttributesPos = { WGPUVertexFormat_Float32x2, sizeof(float) * 0, 0 };
+    WGPUVertexBufferLayout vertexBufferLayouts[] = {
+        makeVertexBufferLayout(&vertexAttributesPos, 1, sizeof(float) * 2)
+    };
+
+    // bind groups
+    WGPUBindGroupLayout bindGroupLayouts[] = {
+        WgBindGroupCanvas::getLayout(device),
+        WgBindGroupPaint::getLayout(device)
+    };
+
+    // stencil function
+    WGPUCompareFunction stencilFunction = WGPUCompareFunction_NotEqual;
+    WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero;
+
+    // shader source and labels
+    auto shaderSource = cShaderSource_PipelineFill;
+    auto shaderLabel = "The shader fill";
+    auto pipelineLabel = "The render pipeline clip mask";
+
+    // allocate all pipeline handles
+    allocate(device, WgPipelineBlendType::SrcOver, WGPUColorWriteMask_All,
+             vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts),
+             bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts),
+             stencilFunction, stencilOperation, stencilFunction, stencilOperation,
+             shaderSource, shaderLabel, pipelineLabel);
+}
+
+
 void WgPipelineSolid::initialize(WGPUDevice device, WgPipelineBlendType blendType)
 {
     // vertex and buffers settings
@@ -148,7 +180,7 @@ void WgPipelineSolid::initialize(WGPUDevice device, WgPipelineBlendType blendTyp
     WGPUCompareFunction stencilFunction = WGPUCompareFunction_NotEqual;
     WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero;
 
-    // sheder source and labels
+    // shader source and labels
     auto shaderSource = cShaderSource_PipelineSolid;
     auto shaderLabel = "The shader solid color";
     auto pipelineLabel = "The render pipeline solid color";
@@ -181,7 +213,7 @@ void WgPipelineLinear::initialize(WGPUDevice device, WgPipelineBlendType blendTy
     WGPUCompareFunction stencilFunction = WGPUCompareFunction_NotEqual;
     WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero;
 
-    // sheder source and labels
+    // shader source and labels
     auto shaderSource = cShaderSource_PipelineLinear;
     auto shaderLabel = "The shader linear gradient";
     auto pipelineLabel = "The render pipeline linear gradient";
@@ -214,7 +246,7 @@ void WgPipelineRadial::initialize(WGPUDevice device, WgPipelineBlendType blendTy
     WGPUCompareFunction stencilFunction = WGPUCompareFunction_NotEqual;
     WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero;
 
-    // sheder source and labels
+    // shader source and labels
     auto shaderSource = cShaderSource_PipelineRadial;
     auto shaderLabel = "The shader radial gradient";
     auto pipelineLabel = "The render pipeline radial gradient";
@@ -249,7 +281,7 @@ void WgPipelineImage::initialize(WGPUDevice device, WgPipelineBlendType blendTyp
     WGPUCompareFunction stencilFunction = WGPUCompareFunction_Always;
     WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero;
 
-    // sheder source and labels
+    // shader source and labels
     auto shaderSource = cShaderSource_PipelineImage;
     auto shaderLabel = "The shader image";
     auto pipelineLabel = "The render pipeline image";
@@ -273,7 +305,7 @@ void WgPipelineClear::initialize(WGPUDevice device)
         WgBindGroupTextureStorageRgba::getLayout(device)
     };
 
-    // sheder source and labels
+    // shader source and labels
     auto shaderSource = cShaderSource_PipelineComputeClear;
     auto shaderLabel = "The compute shader clear";
     auto pipelineLabel = "The compute pipeline clear";
@@ -295,7 +327,7 @@ void WgPipelineBlend::initialize(WGPUDevice device, const char *shaderSource)
         WgBindGroupOpacity::getLayout(device)
     };
 
-    // sheder source and labels
+    // shader source and labels
     auto shaderLabel = "The compute shader blend";
     auto pipelineLabel = "The compute pipeline blend";
 
@@ -306,7 +338,46 @@ void WgPipelineBlend::initialize(WGPUDevice device, const char *shaderSource)
 }
 
 
-void WgPipelineComposeBlend::initialize(WGPUDevice device)
+void WgPipelineBlendMask::initialize(WGPUDevice device, const char *shaderSource)
+{
+    // bind groups and layouts
+    WGPUBindGroupLayout bindGroupLayouts[] = {
+        WgBindGroupTexComposeBlend::getLayout(device),
+        WgBindGroupBlendMethod::getLayout(device),
+        WgBindGroupOpacity::getLayout(device)
+    };
+
+    // shader source and labels
+    auto shaderLabel = "The compute shader blend mask";
+    auto pipelineLabel = "The compute pipeline blend mask";
+
+    // allocate all pipeline handles
+    allocate(device,
+             bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts),
+             shaderSource, shaderLabel, pipelineLabel);
+}
+
+
+void WgPipelineMaskCompose::initialize(WGPUDevice device)
+{
+    // bind groups and layouts
+    WGPUBindGroupLayout bindGroupLayouts[] = {
+        WgBindGroupTexMaskCompose::getLayout(device),
+    };
+
+    // shader source and labels
+    auto shaderSource = cShaderSource_PipelineComputeMaskCompose;
+    auto shaderLabel = "The compute shader mask compose";
+    auto pipelineLabel = "The compute pipeline mask compose";
+
+    // allocate all pipeline handles
+    allocate(device,
+             bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts),
+             shaderSource, shaderLabel, pipelineLabel);
+}
+
+
+void WgPipelineCompose::initialize(WGPUDevice device)
 {
     // bind groups and layouts
     WGPUBindGroupLayout bindGroupLayouts[] = {
@@ -316,8 +387,8 @@ void WgPipelineComposeBlend::initialize(WGPUDevice device)
         WgBindGroupOpacity::getLayout(device)
     };
 
-    // sheder source and labels
+    // shader source and labels
-    auto shaderSource = cShaderSource_PipelineComputeComposeBlend;
+    auto shaderSource = cShaderSource_PipelineComputeCompose;
     auto shaderLabel = "The compute shader compose blend";
     auto pipelineLabel = "The compute pipeline compose blend";
 
@@ -336,7 +407,7 @@ void WgPipelineAntiAliasing::initialize(WGPUDevice device)
         WgBindGroupTextureStorageBgra::getLayout(device)
     };
 
-    // sheder source and labels
+    // shader source and labels
     auto shaderSource = cShaderSource_PipelineComputeAntiAlias;
     auto shaderLabel = "The compute shader anti-aliasing";
     auto pipelineLabel = "The compute pipeline anti-aliasing";
@@ -357,6 +428,7 @@ void WgPipelines::initialize(WgContext& context)
     fillShapeWinding.initialize(context.device);
     fillShapeEvenOdd.initialize(context.device);
     fillStroke.initialize(context.device);
+    clipMask.initialize(context.device);
     for (uint8_t type = (uint8_t)WgPipelineBlendType::SrcOver; type <= (uint8_t)WgPipelineBlendType::Custom; type++) {
         solid[type].initialize(context.device, (WgPipelineBlendType)type);
         linear[type].initialize(context.device, (WgPipelineBlendType)type);
@@ -368,7 +440,11 @@ void WgPipelines::initialize(WgContext& context)
     computeBlendSolid.initialize(context.device, cShaderSource_PipelineComputeBlendSolid);
     computeBlendGradient.initialize(context.device, cShaderSource_PipelineComputeBlendGradient);
     computeBlendImage.initialize(context.device, cShaderSource_PipelineComputeBlendImage);
-    computeComposeBlend.initialize(context.device);
+    computeBlendSolidMask.initialize(context.device, cShaderSource_PipelineComputeBlendSolidMask);
+    computeBlendGradientMask.initialize(context.device, cShaderSource_PipelineComputeBlendGradientMask);
+    computeBlendImageMask.initialize(context.device, cShaderSource_PipelineComputeBlendImageMask);
+    computeMaskCompose.initialize(context.device);
+    computeCompose.initialize(context.device);
     computeAntiAliasing.initialize(context.device);
     // store pipelines to context
     context.pipelines = this;
@@ -377,7 +453,8 @@ void WgPipelines::initialize(WgContext& context)
 
 void WgPipelines::release()
 {
-    WgBindGroupTexComposeBlend::layout = nullptr;
+    WgBindGroupTexMaskCompose::releaseLayout();
+    WgBindGroupTexComposeBlend::releaseLayout();
     WgBindGroupTextureSampled::releaseLayout();
     WgBindGroupTextureStorageBgra::releaseLayout();
     WgBindGroupTextureStorageRgba::releaseLayout();
@@ -391,7 +468,11 @@ void WgPipelines::release()
     WgBindGroupCanvas::releaseLayout();
     // compute pipelines
     computeAntiAliasing.release();
-    computeComposeBlend.release();
+    computeCompose.release();
+    computeMaskCompose.release();
+    computeBlendImageMask.release();
+    computeBlendGradientMask.release();
+    computeBlendSolidMask.release();
     computeBlendImage.release();
     computeBlendGradient.release();
     computeBlendSolid.release();
@@ -403,6 +484,7 @@ void WgPipelines::release()
         linear[type].release();
         solid[type].release();
     }
+    clipMask.release();
     fillStroke.release();
     fillShapeEvenOdd.release();
     fillShapeWinding.release();

src/renderer/wg_engine/tvgWgPipelines.h

@@ -62,6 +62,17 @@ struct WgPipelineFillStroke: public WgRenderPipeline
     }
 };
 
+struct WgPipelineClipMask: public WgRenderPipeline
+{
+    void initialize(WGPUDevice device) override;
+    void use(WGPURenderPassEncoder encoder, WgBindGroupCanvas& groupCanvas,WgBindGroupPaint& groupPaint)
+    {
+        set(encoder);
+        groupCanvas.set(encoder, 0);
+        groupPaint.set(encoder, 1);
+    }
+};
+
 struct WgPipelineSolid: public WgRenderPipeline
 {
     void initialize(WGPUDevice device) override {}
@@ -144,7 +155,32 @@ struct WgPipelineBlend: public WgComputePipeline
 };
 
 
-struct WgPipelineComposeBlend: public WgComputePipeline
+struct WgPipelineBlendMask: public WgComputePipeline
+{
+    void initialize(WGPUDevice device) override { assert(false); };
+    void initialize(WGPUDevice device, const char *shaderSource);
+    void use(WGPUComputePassEncoder encoder, WgBindGroupTexComposeBlend& groupTexs, WgBindGroupBlendMethod& blendMethod, WgBindGroupOpacity& groupOpacity)
+    {
+        set(encoder);
+        groupTexs.set(encoder, 0);
+        blendMethod.set(encoder, 1);
+        groupOpacity.set(encoder, 2);
+    }
+};
+
+
+struct WgPipelineMaskCompose: public WgComputePipeline
+{
+    void initialize(WGPUDevice device) override;
+    void use(WGPUComputePassEncoder encoder, WgBindGroupTexMaskCompose& groupTexs)
+    {
+        set(encoder);
+        groupTexs.set(encoder, 0);
+    }
+};
+
+
+struct WgPipelineCompose: public WgComputePipeline
 {
     void initialize(WGPUDevice device) override;
     void use(WGPUComputePassEncoder encoder, WgBindGroupTexComposeBlend& groupTexs, WgBindGroupCompositeMethod& groupComposeMethod, WgBindGroupBlendMethod& groupBlendMethod, WgBindGroupOpacity& groupOpacity)
@@ -180,6 +216,7 @@ struct WgPipelines
     WgPipelineFillShapeEvenOdd fillShapeEvenOdd;
     WgPipelineFillStroke fillStroke;
     // fill pipelines
+    WgPipelineClipMask clipMask;
     WgPipelineSolid solid[3];
     WgPipelineLinear linear[3];
     WgPipelineRadial radial[3];
@@ -189,7 +226,11 @@ struct WgPipelines
     WgPipelineBlend computeBlendSolid;
     WgPipelineBlend computeBlendGradient;
     WgPipelineBlend computeBlendImage;
-    WgPipelineComposeBlend computeComposeBlend;
+    WgPipelineBlendMask computeBlendSolidMask;
+    WgPipelineBlendMask computeBlendGradientMask;
+    WgPipelineBlendMask computeBlendImageMask;
+    WgPipelineMaskCompose computeMaskCompose;
+    WgPipelineCompose computeCompose;
     WgPipelineAntiAliasing computeAntiAliasing;
 
     void initialize(WgContext& context);

src/renderer/wg_engine/tvgWgRenderData.cpp

@@ -262,8 +262,17 @@ void WgRenderSettings::release(WgContext& context)
 void WgRenderDataPaint::release(WgContext& context)
 {
     bindGroupPaint.release();
+    clips.clear();
 };
 
+
+void WgRenderDataPaint::updateClips(tvg::Array<tvg::RenderData> &clips) {
+    this->clips.clear();
+    for (uint32_t i = 0; i < clips.count; i++)
+        if (clips[i])
+            this->clips.push((WgRenderDataPaint*)clips[i]);
+}
+
 //***********************************************************************
 // WgRenderDataShape
 //***********************************************************************
@@ -371,6 +380,7 @@ void WgRenderDataShape::releaseMeshes(WgContext &context)
     meshGroupShapes.release(context);
     pMin = {FLT_MAX, FLT_MAX};
     pMax = {0.0f, 0.0f};
+    clips.clear();
 }
 
 
@@ -407,6 +417,7 @@ void WgRenderDataShapePool::free(WgContext& context, WgRenderDataShape* dataShap
     dataShape->meshGroupShapesBBox.release(context);
     dataShape->meshGroupStrokes.release(context);
     dataShape->meshGroupStrokesBBox.release(context);
+    dataShape->clips.clear();
     mPool.push(dataShape);
 }
 

src/renderer/wg_engine/tvgWgRenderData.h

@@ -88,10 +88,13 @@ struct WgRenderDataPaint
     WgBindGroupPaint bindGroupPaint{};
     RenderRegion viewport{};
     float opacity{};
+    Array<WgRenderDataPaint*> clips;
 
     virtual ~WgRenderDataPaint() {};
     virtual void release(WgContext& context);
     virtual Type type() { return Type::Undefined; };
+
+    void updateClips(tvg::Array<tvg::RenderData> &clips);
 };
 
 struct WgRenderDataShape: public WgRenderDataPaint

src/renderer/wg_engine/tvgWgRenderTarget.cpp

@@ -85,7 +85,6 @@ void WgRenderStorage::renderShape(WgContext& context, WgRenderDataShape* renderD
 {
     assert(renderData);
     assert(mRenderPassEncoder);
-    // draw strokes
     if (renderData->strokeFirst) {
         drawStroke(context, renderData, blendType);
         drawShape(context, renderData, blendType);
@@ -102,12 +101,21 @@ void WgRenderStorage::renderPicture(WgContext& context, WgRenderDataPicture* ren
     assert(mRenderPassEncoder);
     uint8_t blend = (uint8_t)blendType;
     auto& vp = renderData->viewport;
+    if ((vp.w <= 0) || (vp.h <= 0)) return;
     wgpuRenderPassEncoderSetScissorRect(mRenderPassEncoder, vp.x * samples, vp.y * samples, vp.w * samples, vp.h * samples);
     wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 0);
     mPipelines->image[blend].use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint, renderData->bindGroupPicture);
     renderData->meshData.drawImage(context, mRenderPassEncoder);
 }
 
+void WgRenderStorage::renderClipPath(WgContext& context, WgRenderDataPaint* renderData) {
+    assert(renderData);
+    assert(mRenderPassEncoder);
+    if (renderData->type() == Type::Shape)
+        drawShapeClipPath(context, (WgRenderDataShape*)renderData);
+    else if (renderData->type() == Type::Picture)
+        drawPictureClipPath(context, (WgRenderDataPicture*)renderData);
+}
 
 void WgRenderStorage::drawShape(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType)
 {
@@ -115,10 +123,12 @@ void WgRenderStorage::drawShape(WgContext& context, WgRenderDataShape* renderDat
     assert(mRenderPassEncoder);
     assert(renderData->meshGroupShapes.meshes.count == renderData->meshGroupShapesBBox.meshes.count);
     if (renderData->renderSettingsShape.skip) return;
+    // apply viewport
     auto& vp = renderData->viewport;
+    if ((vp.w <= 0) || (vp.h <= 0)) return;
     wgpuRenderPassEncoderSetScissorRect(mRenderPassEncoder, vp.x * samples, vp.y * samples, vp.w * samples, vp.h * samples);
-    wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 0);
     // setup fill rule
+    wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 0);
     if (renderData->fillRule == FillRule::Winding)
         mPipelines->fillShapeWinding.use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint);
     else
@@ -145,7 +155,9 @@ void WgRenderStorage::drawStroke(WgContext& context, WgRenderDataShape* renderDa
     assert(mRenderPassEncoder);
     assert(renderData->meshGroupStrokes.meshes.count == renderData->meshGroupStrokesBBox.meshes.count);
     if (renderData->renderSettingsStroke.skip) return;
+    // apply viewport
     auto& vp = renderData->viewport;
+    if ((vp.w <= 0) || (vp.h <= 0)) return;
     wgpuRenderPassEncoderSetScissorRect(mRenderPassEncoder, vp.x * samples, vp.y * samples, vp.w * samples, vp.h * samples);
     // draw stroke geometry
     uint8_t blend = (uint8_t)blendType;
@@ -169,6 +181,35 @@ void WgRenderStorage::drawStroke(WgContext& context, WgRenderDataShape* renderDa
 }
 
 
+void WgRenderStorage::drawShapeClipPath(WgContext& context, WgRenderDataShape* renderData) {
+    assert(renderData);
+    assert(renderData->type() == Type::Shape);
+    assert(renderData->meshGroupShapes.meshes.count == renderData->meshGroupShapesBBox.meshes.count);
+    // draw shape geometry
+    wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 0);
+    // setup fill rule
+    if (renderData->fillRule == FillRule::Winding)
+        mPipelines->fillShapeWinding.use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint);
+    else
+        mPipelines->fillShapeEvenOdd.use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint);
+    // draw to stencil (first pass)
+    for (uint32_t i = 0; i < renderData->meshGroupShapes.meshes.count; i++)
+        renderData->meshGroupShapes.meshes[i]->drawFan(context, mRenderPassEncoder);
+    // fill shape geometry (second pass)
+    mPipelines->clipMask.use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint);
+    renderData->meshDataBBox.drawFan(context, mRenderPassEncoder);
+}
+
+void WgRenderStorage::drawPictureClipPath(WgContext& context, WgRenderDataPicture* renderData) {
+    assert(renderData);
+    assert(renderData->type() == Type::Picture);
+    assert(mRenderPassEncoder);
+    wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 0);
+    mPipelines->clipMask.use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint);
+    renderData->meshData.drawImage(context, mRenderPassEncoder);
+}
+
+
 void WgRenderStorage::clear(WGPUCommandEncoder commandEncoder)
 {
     assert(commandEncoder);
@@ -179,7 +220,14 @@ void WgRenderStorage::clear(WGPUCommandEncoder commandEncoder)
 }
 
 
-void WgRenderStorage::blend(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc, WgPipelineBlend* pipeline, WgBindGroupBlendMethod* blendMethod, WgBindGroupOpacity* opacity) {
+void WgRenderStorage::blend(
+    WgContext& context,
+    WGPUCommandEncoder commandEncoder,
+    WgPipelineBlend* pipeline,
+    WgRenderStorage* targetSrc,
+    WgBindGroupBlendMethod* blendMethod,
+    WgBindGroupOpacity* opacity)
+{
     assert(commandEncoder);
     assert(targetSrc);
     assert(pipeline);
@@ -190,7 +238,46 @@ void WgRenderStorage::blend(WGPUCommandEncoder commandEncoder, WgRenderStorage*
 }
 
 
-void WgRenderStorage::composeBlend(
+void WgRenderStorage::blendMask(
+    WgContext& context,
+    WGPUCommandEncoder commandEncoder,
+    WgPipelineBlendMask* pipeline,
+    WgRenderStorage* texMsk,
+    WgRenderStorage* texSrc,
+    WgBindGroupBlendMethod* blendMethod,
+    WgBindGroupOpacity* opacity)
+{
+    assert(commandEncoder);
+    assert(texSrc);
+    assert(texMsk);
+    WgBindGroupTexComposeBlend composeBlend;
+    composeBlend.initialize(context.device, context.queue, texSrc->texViewColor, texMsk->texViewColor, texViewColor);
+    WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder);
+    pipeline->use(computePassEncoder, composeBlend, *blendMethod, *opacity);
+    dispatchWorkgroups(computePassEncoder);
+    endComputePass(computePassEncoder);
+    composeBlend.release();
+};
+
+
+void WgRenderStorage::maskCompose(
+        WgContext& context,
+        WGPUCommandEncoder commandEncoder,
+        WgRenderStorage* texMsk0)
+{
+    assert(commandEncoder);
+    assert(texMsk0);
+    WgBindGroupTexMaskCompose maskCompose;
+    maskCompose.initialize(context.device, context.queue, texMsk0->texViewColor, texViewColor);
+    WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder);
+    mPipelines->computeMaskCompose.use(computePassEncoder, maskCompose);
+    dispatchWorkgroups(computePassEncoder);
+    endComputePass(computePassEncoder);
+    maskCompose.release();
+}
+
+
+void WgRenderStorage::compose(
     WgContext& context,
     WGPUCommandEncoder commandEncoder,
     WgRenderStorage* texSrc,
@@ -205,7 +292,7 @@ void WgRenderStorage::composeBlend(
     WgBindGroupTexComposeBlend composeBlend;
     composeBlend.initialize(context.device, context.queue, texSrc->texViewColor, texMsk->texViewColor, texViewColor);
     WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder);
-    mPipelines->computeComposeBlend.use(computePassEncoder, composeBlend, *composeMethod, *blendMethod, *opacity);
+    mPipelines->computeCompose.use(computePassEncoder, composeBlend, *composeMethod, *blendMethod, *opacity);
     dispatchWorkgroups(computePassEncoder);
     endComputePass(computePassEncoder);
     composeBlend.release();

src/renderer/wg_engine/tvgWgRenderTarget.h

@@ -52,10 +52,29 @@ public:
 
     void renderShape(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType);
     void renderPicture(WgContext& context, WgRenderDataPicture* renderData, WgPipelineBlendType blendType);
+    void renderClipPath(WgContext& context, WgRenderDataPaint* renderData);
 
     void clear(WGPUCommandEncoder commandEncoder);
-    void blend(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc, WgPipelineBlend* pipeline, WgBindGroupBlendMethod* blendMethod, WgBindGroupOpacity* opacity);
+    void blend(
-    void composeBlend(
+        WgContext& context,
+        WGPUCommandEncoder commandEncoder,
+        WgPipelineBlend* pipeline,
+        WgRenderStorage* targetSrc,
+        WgBindGroupBlendMethod* blendMethod,
+        WgBindGroupOpacity* opacity);
+    void blendMask(
+        WgContext& context,
+        WGPUCommandEncoder commandEncoder,
+        WgPipelineBlendMask* pipeline,
+        WgRenderStorage* texMsk,
+        WgRenderStorage* texSrc,
+        WgBindGroupBlendMethod* blendMethod,
+        WgBindGroupOpacity* opacity);
+    void maskCompose(
+        WgContext& context,
+        WGPUCommandEncoder commandEncoder,
+        WgRenderStorage* texMsk0);
+    void compose(
         WgContext& context,
         WGPUCommandEncoder commandEncoder,
         WgRenderStorage* texMsk,
@@ -68,6 +87,9 @@ private:
     void drawShape(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType);
     void drawStroke(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType);
 
+    void drawShapeClipPath(WgContext& context, WgRenderDataShape* renderData);
+    void drawPictureClipPath(WgContext& context, WgRenderDataPicture* renderData);
+
     void dispatchWorkgroups(WGPUComputePassEncoder computePassEncoder);
 
     WGPUComputePassEncoder beginComputePass(WGPUCommandEncoder commandEncoder);

src/renderer/wg_engine/tvgWgRenderer.cpp

@@ -67,6 +67,7 @@ void WgRenderer::release()
     mBlendMethodPool.release(mContext);
     mOpacityPool.release(mContext);
     mRenderStorageRoot.release(mContext);
+    mRenderStorageMask.release(mContext);
     mRenderStorageScreen.release(mContext);
     mRenderStorageInterm.release(mContext);
     mPipelines.release();
@@ -119,6 +120,9 @@ RenderData WgRenderer::prepare(const RenderShape& rshape, RenderData data, const
     if (rshape.stroke)
         renderDataShape->renderSettingsStroke.update(mContext, rshape.stroke->fill, rshape.stroke->color, flags);
 
+    // store clips data
+    renderDataShape->updateClips(clips);
+
     return renderDataShape;
 }
 
@@ -159,6 +163,9 @@ RenderData WgRenderer::prepare(Surface* surface, const RenderMesh* mesh, RenderD
             renderDataPicture->imageData.textureView);
     }
 
+    // store clips data
+    renderDataPicture->updateClips(clips);
+
     return renderDataPicture;
 }
 
@@ -175,6 +182,19 @@ bool WgRenderer::preRender()
 }
 
 
+void WgRenderer::renderClipPath(Array<WgRenderDataPaint*>& clips)
+{
+    for (uint32_t i = 0; i < clips.count; i++) {
+        renderClipPath(clips[i]->clips);
+        // render image to render target
+        mRenderStorageInterm.beginRenderPass(mCommandEncoder, true);
+        mRenderStorageInterm.renderClipPath(mContext, clips[i]);
+        mRenderStorageInterm.endRenderPass();
+        mRenderStorageMask.maskCompose(mContext, mCommandEncoder, &mRenderStorageInterm);
+    }
+}
+
+
 bool WgRenderer::renderShape(RenderData data)
 {
     // get current render storage
@@ -183,11 +203,34 @@ bool WgRenderer::renderShape(RenderData data)
     WgPipelineBlendType blendType = WgPipelines::blendMethodToBlendType(mBlendMethod);
     WgRenderStorage* renderStorage = mRenderStorageStack.last();
     assert(renderStorage);
+    // use masked blend
+    if (dataShape->clips.count > 0) {
+        // terminate current render pass
+        renderStorage->endRenderPass();
+        // render clip path
+        mRenderStorageMask.beginRenderPass(mCommandEncoder, true);
+        mRenderStorageMask.renderClipPath(mContext, dataShape->clips[0]);
+        mRenderStorageMask.endRenderPass();
+        renderClipPath(dataShape->clips);
+        // render image to render target
+        mRenderStorageInterm.beginRenderPass(mCommandEncoder, true);
+        mRenderStorageInterm.renderShape(mContext, dataShape, blendType);
+        mRenderStorageInterm.endRenderPass();
+        // blend shape with current render storage
+        WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod);
+        WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, 255);
+        WgPipelineBlendMask* pipeline = &mContext.pipelines->computeBlendSolidMask;
+        if (dataShape->renderSettingsShape.fillType != WgRenderSettingsType::Solid)
+            pipeline = &mContext.pipelines->computeBlendGradientMask;
+        renderStorage->blendMask(mContext, mCommandEncoder, 
+            pipeline, &mRenderStorageMask, &mRenderStorageInterm, blendMethod, opacity);
+        // restore current render pass
+        renderStorage->beginRenderPass(mCommandEncoder, false);
     // use hardware blend
-    if (WgPipelines::isBlendMethodSupportsHW(mBlendMethod))
+    } else if (WgPipelines::isBlendMethodSupportsHW(mBlendMethod)) {
         renderStorage->renderShape(mContext, dataShape, blendType);
     // use custom blend
-    else {
+    } else {
         // terminate current render pass
         renderStorage->endRenderPass();
         // render image to render target
@@ -200,7 +243,8 @@ bool WgRenderer::renderShape(RenderData data)
         WgPipelineBlend* pipeline = &mContext.pipelines->computeBlendSolid;
         if (dataShape->renderSettingsShape.fillType != WgRenderSettingsType::Solid)
             pipeline = &mContext.pipelines->computeBlendGradient;
-        renderStorage->blend(mCommandEncoder, &mRenderStorageInterm, pipeline, blendMethod, opacity);
+        renderStorage->blend(mContext, mCommandEncoder, 
+            pipeline, &mRenderStorageInterm, blendMethod, opacity);
         // restore current render pass
         renderStorage->beginRenderPass(mCommandEncoder, false);
     }
@@ -210,26 +254,50 @@ bool WgRenderer::renderShape(RenderData data)
 
 bool WgRenderer::renderImage(RenderData data)
 {
+    // get current render storage
+    WgRenderDataPicture *dataPicture = (WgRenderDataPicture*)data;
     // get current render storage
     WgPipelineBlendType blendType = WgPipelines::blendMethodToBlendType(mBlendMethod);
     WgRenderStorage* renderStorage = mRenderStorageStack.last();
     assert(renderStorage);
+        // use masked blend
+    if (dataPicture->clips.count > 0) {
+        // terminate current render pass
+        renderStorage->endRenderPass();
+        // render clip path
+        mRenderStorageMask.beginRenderPass(mCommandEncoder, true);
+        mRenderStorageMask.renderClipPath(mContext, dataPicture->clips[0]);
+        mRenderStorageMask.endRenderPass();
+        renderClipPath(dataPicture->clips);
+        // render image to render target
+        mRenderStorageInterm.beginRenderPass(mCommandEncoder, true);
+        mRenderStorageInterm.renderPicture(mContext, dataPicture, blendType);
+        mRenderStorageInterm.endRenderPass();
+        // blend shape with current render storage
+        WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod);
+        WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, 255);
+        WgPipelineBlendMask* pipeline = &mContext.pipelines->computeBlendImageMask;
+        renderStorage->blendMask(mContext, mCommandEncoder, 
+            pipeline, &mRenderStorageMask, &mRenderStorageInterm, blendMethod, opacity);
+        // restore current render pass
+        renderStorage->beginRenderPass(mCommandEncoder, false);
     // use hardware blend
-    if (WgPipelines::isBlendMethodSupportsHW(mBlendMethod))
+    } else if (WgPipelines::isBlendMethodSupportsHW(mBlendMethod))
-        renderStorage->renderPicture(mContext, (WgRenderDataPicture *)data, blendType);
+        renderStorage->renderPicture(mContext, dataPicture, blendType);
     // use custom blend
     else {
         // terminate current render pass
         renderStorage->endRenderPass();
         // render image to render target
         mRenderStorageInterm.beginRenderPass(mCommandEncoder, true);
-        mRenderStorageInterm.renderPicture(mContext, (WgRenderDataPicture *)data, blendType);
+        mRenderStorageInterm.renderPicture(mContext, dataPicture, blendType);
         mRenderStorageInterm.endRenderPass();
         // blend shape with current render storage
         WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod);
         WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, 255);
         WgPipelineBlend* pipeline = &mContext.pipelines->computeBlendImage;
-        renderStorage->blend(mCommandEncoder, &mRenderStorageInterm, pipeline, blendMethod, opacity);
+        renderStorage->blend(mContext, mCommandEncoder, 
+            pipeline, &mRenderStorageInterm, blendMethod, opacity);
         // restore current render pass
         renderStorage->beginRenderPass(mCommandEncoder, false);
     }
@@ -258,7 +326,7 @@ void WgRenderer::dispose(RenderData data) {
 
 RenderRegion WgRenderer::region(TVG_UNUSED RenderData data)
 {
-    return { 0, 0, INT32_MAX, INT32_MAX };
+    return { 0, 0, (int32_t)mTargetSurface.w, (int32_t)mTargetSurface.h };
 }
 
 
@@ -355,6 +423,7 @@ bool WgRenderer::target(WGPUInstance instance, WGPUSurface surface, uint32_t w,
 
     initialize();
     mRenderStorageInterm.initialize(mContext, w, h, WG_SSAA_SAMPLES);
+    mRenderStorageMask.initialize(mContext, w, h, WG_SSAA_SAMPLES);
     mRenderStorageRoot.initialize(mContext, w, h, WG_SSAA_SAMPLES);
     mRenderStorageScreen.initialize(mContext, w, h, 1, WGPUTextureFormat_BGRA8Unorm);
     return true;
@@ -400,7 +469,9 @@ bool WgRenderer::endComposite(TVG_UNUSED Compositor* cmp)
         // blent scene to current render storage
         WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, comp->blendMethod);
         WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, comp->opacity);
-        mRenderStorageStack.last()->blend(mCommandEncoder, renderStorageSrc, &mContext.pipelines->computeBlendImage, blendMethod, opacity);
+        mRenderStorageStack.last()->blend(mContext, mCommandEncoder,
+            &mContext.pipelines->computeBlendImage,
+            renderStorageSrc, blendMethod, opacity);
 
         // back render targets to the pool
         mRenderStoragePool.free(mContext, renderStorageSrc);
@@ -424,7 +495,7 @@ bool WgRenderer::endComposite(TVG_UNUSED Compositor* cmp)
 
         // compose and blend
         // dest = blend(dest, compose(src, msk, composeMethod), blendMethod, opacity)
-        renderStorageDst->composeBlend(mContext, mCommandEncoder,
+        renderStorageDst->compose(mContext, mCommandEncoder,
             renderStorageSrc, renderStorageMsk,
             composeMethod, blendMethod, opacity);
 

src/renderer/wg_engine/tvgWgRenderer.h

@@ -62,12 +62,14 @@ private:
     void initialize();
     void release();
     void clearDisposes();
+    void renderClipPath(Array<WgRenderDataPaint*>& clips);
 
-    WGPUCommandEncoder mCommandEncoder{};    // render handles
+    WGPUCommandEncoder mCommandEncoder{};
-    WgRenderStorage mRenderStorageInterm;    // intermidiate buffer to render
+    WgRenderStorage mRenderStorageInterm; // intermidiate buffer to render
-    WgRenderStorage mRenderStorageRoot;    // root render storage
+    WgRenderStorage mRenderStorageMask;   // buffer to render mask 
-    WgRenderStorage mRenderStorageScreen;    // storage with data after antializing
+    WgRenderStorage mRenderStorageRoot;   // root render storage
-    WgRenderStoragePool mRenderStoragePool;    // pool to hold render tree storages
+    WgRenderStorage mRenderStorageScreen; // storage with data after antializing
+    WgRenderStoragePool mRenderStoragePool; // pool to hold render tree storages
     WgBindGroupOpacityPool mOpacityPool;    // opacity, blend methods and composite methods pool
     WgBindGroupBlendMethodPool mBlendMethodPool;
     WgBindGroupCompositeMethodPool mCompositeMethodPool;

src/renderer/wg_engine/tvgWgShaderSrc.cpp

@@ -53,8 +53,8 @@ fn vs_main(in: VertexInput) -> VertexOutput {
 }
 
 @fragment
-fn fs_main(in: VertexOutput) -> void {
+fn fs_main(in: VertexOutput) -> @location(0) vec4f {
-    // nothing to draw, just stencil value
+    return vec4f(0.0, 0.0, 0.0, 1.0);
 }
 );
 
@@ -393,6 +393,34 @@ fn cs_main( @builtin(global_invocation_id) id: vec3u) {
     var Ra: f32    = 1.0;
 );
 
+const std::string strBlendMaskShaderHeader = WG_SHADER_SOURCE(
+@group(0) @binding(0) var imageSrc : texture_storage_2d<rgba8unorm, read>;
+@group(0) @binding(1) var imageMsk : texture_storage_2d<rgba8unorm, read>;
+@group(0) @binding(2) var imageDst : texture_storage_2d<rgba8unorm, read_write>;
+@group(1) @binding(0) var<uniform> blendMethod : u32;
+@group(2) @binding(0) var<uniform> opacity : f32;
+
+@compute @workgroup_size(8, 8)
+fn cs_main( @builtin(global_invocation_id) id: vec3u) {
+    let texSize = textureDimensions(imageSrc);
+    if ((id.x >= texSize.x) || (id.y >= texSize.y)) { return; };
+
+    let colorMsk = textureLoad(imageMsk, id.xy);
+    if (colorMsk.a == 0.0) { return; }
+    let colorSrc = textureLoad(imageSrc, id.xy);
+    if (colorSrc.a == 0.0) { return; }
+    let colorDst = textureLoad(imageDst, id.xy);
+
+    var One: vec3f = vec3(1.0);
+    var So: f32    = opacity;
+    var Sc: vec3f  = colorSrc.rgb;
+    var Sa: f32    = colorSrc.a;
+    var Dc: vec3f  = colorDst.rgb;
+    var Da: f32    = colorDst.a;
+    var Rc: vec3f  = colorDst.rgb;
+    var Ra: f32    = 1.0;
+);
+
 const std::string strBlendShaderPreConditionsGradient = WG_SHADER_SOURCE(
     Sc = Sc + Dc.rgb * (1.0 - Sa);
     Sa = Sa + Da     * (1.0 - Sa);
@@ -484,6 +512,31 @@ const std::string strComputeBlendImage =
     strBlendShaderFooter;
 const char* cShaderSource_PipelineComputeBlendImage = strComputeBlendImage.c_str();
 
+// pipeline shader modules blend solid mask
+const std::string strComputeBlendSolidMask =
+    strBlendMaskShaderHeader + 
+    strBlendShaderBlendMethod + 
+    strBlendShaderFooter;
+const char* cShaderSource_PipelineComputeBlendSolidMask = strComputeBlendSolidMask.c_str();
+
+// pipeline shader modules blend gradient mask
+const std::string strComputeBlendGradientMask =
+    strBlendMaskShaderHeader + 
+    strBlendShaderPreConditionsGradient + 
+    strBlendShaderBlendMethod + 
+    strBlendShaderPostConditionsGradient + 
+    strBlendShaderFooter;
+const char* cShaderSource_PipelineComputeBlendGradientMask = strComputeBlendGradientMask.c_str();
+
+// pipeline shader modules blend image mask
+const std::string strComputeBlendImageMask =
+    strBlendMaskShaderHeader + 
+    strBlendShaderPreConditionsImage + 
+    strBlendShaderBlendMethod + 
+    strBlendShaderPostConditionsImage + 
+    strBlendShaderFooter;
+const char* cShaderSource_PipelineComputeBlendImageMask = strComputeBlendImageMask.c_str();
+
 // pipeline shader modules clear
 const char* cShaderSource_PipelineComputeClear = WG_SHADER_SOURCE(
 @group(0) @binding(0) var imageDst : texture_storage_2d<rgba8unorm, read_write>;
@@ -494,8 +547,21 @@ fn cs_main( @builtin(global_invocation_id) id: vec3u) {
 }
 );
 
-// pipeline shader modules compose blend
+// pipeline shader modules compose
-const char* cShaderSource_PipelineComputeComposeBlend = WG_SHADER_SOURCE(
+const char* cShaderSource_PipelineComputeMaskCompose = WG_SHADER_SOURCE(
+@group(0) @binding(0) var imageMsk0 : texture_storage_2d<rgba8unorm, read>;
+@group(0) @binding(1) var imageMsk1 : texture_storage_2d<rgba8unorm, read_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(imageMsk1, id.xy, colorMsk0 * colorMsk1);
+}
+);
+
+// pipeline shader modules compose
+const char* cShaderSource_PipelineComputeCompose = WG_SHADER_SOURCE(
 @group(0) @binding(0) var imageSrc : texture_storage_2d<rgba8unorm, read>;
 @group(0) @binding(1) var imageMsk : texture_storage_2d<rgba8unorm, read>;
 @group(0) @binding(2) var imageDst : texture_storage_2d<rgba8unorm, read_write>;

src/renderer/wg_engine/tvgWgShaderSrc.h

@@ -45,7 +45,11 @@ extern const char* cShaderSource_PipelineComputeClear;
 extern const char* cShaderSource_PipelineComputeBlendSolid;
 extern const char* cShaderSource_PipelineComputeBlendGradient;
 extern const char* cShaderSource_PipelineComputeBlendImage;
-extern const char* cShaderSource_PipelineComputeComposeBlend;
+extern const char* cShaderSource_PipelineComputeBlendSolidMask;
+extern const char* cShaderSource_PipelineComputeBlendGradientMask;
+extern const char* cShaderSource_PipelineComputeBlendImageMask;
+extern const char* cShaderSource_PipelineComputeMaskCompose;
+extern const char* cShaderSource_PipelineComputeCompose;
 extern const char* cShaderSource_PipelineComputeAntiAlias;
 
 #endif // _TVG_WG_SHADER_SRC_H_