wg_engine: emscripten and browser support

[issues 2410: emscripten support](https://github.com/thorvg/thorvg/issues/2410)
- Research and prototype the Emscripten build with WebGPU.

meson setup script:
   meson setup builddir -Ddefault_library=static -Dloaders=all -Dsavers="all" -Dbindings="wasm_beta" -Dengines="wg_beta" --cross-file ./cross/wasm_webgpu.txt

src/renderer/wg_engine/tvgWgBindGroups.cpp

@@ -33,11 +33,15 @@ WGPUBindGroupLayout WgBindGroupRadialGradient::layout = nullptr;
 WGPUBindGroupLayout WgBindGroupPicture::layout = nullptr;
 // composition and blending properties groups
 WGPUBindGroupLayout WgBindGroupTexture::layout = nullptr;
-WGPUBindGroupLayout WgBindGroupTextureStorageRgba::layout = nullptr;
+WGPUBindGroupLayout WgBindGroupTextureStorageRgbaRO::layout = nullptr;
-WGPUBindGroupLayout WgBindGroupTextureStorageBgra::layout = nullptr;
+WGPUBindGroupLayout WgBindGroupTextureStorageRgbaWO::layout = nullptr;
+WGPUBindGroupLayout WgBindGroupTextureStorageBgraRO::layout = nullptr;
+WGPUBindGroupLayout WgBindGroupTextureStorageBgraWO::layout = nullptr;
 WGPUBindGroupLayout WgBindGroupTextureSampled::layout = nullptr;
+WGPUBindGroupLayout WgBindGroupTexBlend::layout = nullptr;
+WGPUBindGroupLayout WgBindGroupTexBlendMask::layout = nullptr;
 WGPUBindGroupLayout WgBindGroupTexMaskCompose::layout = nullptr;
-WGPUBindGroupLayout WgBindGroupTexComposeBlend::layout = nullptr;
+WGPUBindGroupLayout WgBindGroupTexCompose::layout = nullptr;
 WGPUBindGroupLayout WgBindGroupOpacity::layout = nullptr;
 WGPUBindGroupLayout WgBindGroupBlendMethod::layout = nullptr;
 WGPUBindGroupLayout WgBindGroupCompositeMethod::layout = nullptr;
@@ -338,14 +342,14 @@ void WgBindGroupTexture::release()
 }
 
 ///////////////////////////////////////////////////////////////////////////////
-// WgBindGroupTextureStorageRgba
+// WgBindGroupTextureStorageRgbaRO
 ///////////////////////////////////////////////////////////////////////////////
 
-WGPUBindGroupLayout WgBindGroupTextureStorageRgba::getLayout(WGPUDevice device)
+WGPUBindGroupLayout WgBindGroupTextureStorageRgbaRO::getLayout(WGPUDevice device)
 {
     if (layout) return layout;
     const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] {
-        makeBindGroupLayoutEntryStorage(0, WGPUStorageTextureAccess_ReadWrite, WGPUTextureFormat_RGBA8Unorm)
+        makeBindGroupLayoutEntryStorage(0, WGPUStorageTextureAccess_ReadOnly, WGPUTextureFormat_RGBA8Unorm)
     };
     layout = createBindGroupLayout(device, bindGroupLayoutEntries, 1);
     assert(layout);
@@ -353,13 +357,13 @@ WGPUBindGroupLayout WgBindGroupTextureStorageRgba::getLayout(WGPUDevice device)
 }
 
 
-void WgBindGroupTextureStorageRgba::releaseLayout()
+void WgBindGroupTextureStorageRgbaRO::releaseLayout()
 {
     releaseBindGroupLayout(layout);
 }
 
 
-void WgBindGroupTextureStorageRgba::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexture)
+void WgBindGroupTextureStorageRgbaRO::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexture)
 {
     release();
     const WGPUBindGroupEntry bindGroupEntries[] {
@@ -370,16 +374,91 @@ void WgBindGroupTextureStorageRgba::initialize(WGPUDevice device, WGPUQueue queu
 }
 
 
-void WgBindGroupTextureStorageRgba::release()
+void WgBindGroupTextureStorageRgbaRO::release()
 {
     releaseBindGroup(mBindGroup);
 }
 
 ///////////////////////////////////////////////////////////////////////////////
-// WgBindGroupTextureStorageBgra
+// WgBindGroupTextureStorageRgbaRO
 ///////////////////////////////////////////////////////////////////////////////
 
-WGPUBindGroupLayout WgBindGroupTextureStorageBgra::getLayout(WGPUDevice device)
+WGPUBindGroupLayout WgBindGroupTextureStorageRgbaWO::getLayout(WGPUDevice device)
+{
+    if (layout) return layout;
+    const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] {
+        makeBindGroupLayoutEntryStorage(0, WGPUStorageTextureAccess_WriteOnly, WGPUTextureFormat_RGBA8Unorm)
+    };
+    layout = createBindGroupLayout(device, bindGroupLayoutEntries, 1);
+    assert(layout);
+    return layout;
+}
+
+
+void WgBindGroupTextureStorageRgbaWO::releaseLayout()
+{
+    releaseBindGroupLayout(layout);
+}
+
+
+void WgBindGroupTextureStorageRgbaWO::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexture)
+{
+    release();
+    const WGPUBindGroupEntry bindGroupEntries[] {
+        makeBindGroupEntryTextureView(0, uTexture)
+    };
+    mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 1);
+    assert(mBindGroup);
+}
+
+void WgBindGroupTextureStorageRgbaWO::release()
+{
+    releaseBindGroup(mBindGroup);
+}
+
+///////////////////////////////////////////////////////////////////////////////
+// WgBindGroupTextureStorageBgraRO
+///////////////////////////////////////////////////////////////////////////////
+
+WGPUBindGroupLayout WgBindGroupTextureStorageBgraRO::getLayout(WGPUDevice device)
+{
+    if (layout) return layout;
+    const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] {
+        makeBindGroupLayoutEntryStorage(0, WGPUStorageTextureAccess_ReadOnly, WGPUTextureFormat_BGRA8Unorm)
+    };
+    layout = createBindGroupLayout(device, bindGroupLayoutEntries, 1);
+    assert(layout);
+    return layout;
+}
+
+
+void WgBindGroupTextureStorageBgraRO::releaseLayout()
+{
+    releaseBindGroupLayout(layout);
+}
+
+
+void WgBindGroupTextureStorageBgraRO::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexture)
+{
+    release();
+    const WGPUBindGroupEntry bindGroupEntries[] {
+        makeBindGroupEntryTextureView(0, uTexture)
+    };
+    mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 1);
+    assert(mBindGroup);
+}
+
+
+void WgBindGroupTextureStorageBgraRO::release()
+{
+    releaseBindGroup(mBindGroup);
+}
+
+///////////////////////////////////////////////////////////////////////////////
+// WgBindGroupTextureStorageBgraWO
+///////////////////////////////////////////////////////////////////////////////
+
+WGPUBindGroupLayout WgBindGroupTextureStorageBgraWO::getLayout(WGPUDevice device)
 {
     if (layout) return layout;
     const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] {
@@ -391,13 +470,13 @@ WGPUBindGroupLayout WgBindGroupTextureStorageBgra::getLayout(WGPUDevice device)
 }
 
 
-void WgBindGroupTextureStorageBgra::releaseLayout()
+void WgBindGroupTextureStorageBgraWO::releaseLayout()
 {
     releaseBindGroupLayout(layout);
 }
 
 
-void WgBindGroupTextureStorageBgra::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexture)
+void WgBindGroupTextureStorageBgraWO::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexture)
 {
     release();
     const WGPUBindGroupEntry bindGroupEntries[] {
@@ -408,7 +487,7 @@ void WgBindGroupTextureStorageBgra::initialize(WGPUDevice device, WGPUQueue queu
 }
 
 
-void WgBindGroupTextureStorageBgra::release()
+void WgBindGroupTextureStorageBgraWO::release()
 {
     releaseBindGroup(mBindGroup);
 }
@@ -454,16 +533,16 @@ void WgBindGroupTextureSampled::release()
 }
 
 ///////////////////////////////////////////////////////////////////////////////
-// WgBindGroupTexComposeBlend
+// WgBindGroupTexBlend
 ///////////////////////////////////////////////////////////////////////////////
 
-WGPUBindGroupLayout WgBindGroupTexComposeBlend::getLayout(WGPUDevice device)
+WGPUBindGroupLayout WgBindGroupTexBlend::getLayout(WGPUDevice device)
 {
     if (layout) return layout;
     const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] {
         makeBindGroupLayoutEntryStorage(0, WGPUStorageTextureAccess_ReadOnly, WGPUTextureFormat_RGBA8Unorm),
         makeBindGroupLayoutEntryStorage(1, WGPUStorageTextureAccess_ReadOnly, WGPUTextureFormat_RGBA8Unorm),
-        makeBindGroupLayoutEntryStorage(2, WGPUStorageTextureAccess_ReadWrite, WGPUTextureFormat_RGBA8Unorm)
+        makeBindGroupLayoutEntryStorage(2, WGPUStorageTextureAccess_WriteOnly, WGPUTextureFormat_RGBA8Unorm)
     };
     layout = createBindGroupLayout(device, bindGroupLayoutEntries, 3);
     assert(layout);
@@ -471,13 +550,13 @@ WGPUBindGroupLayout WgBindGroupTexComposeBlend::getLayout(WGPUDevice device)
 }
 
 
-void WgBindGroupTexComposeBlend::releaseLayout()
+void WgBindGroupTexBlend::releaseLayout()
 {
     releaseBindGroupLayout(layout);
 }
 
 
-void WgBindGroupTexComposeBlend::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexSrc, WGPUTextureView uTexMsk, WGPUTextureView uTexDst)
+void WgBindGroupTexBlend::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexSrc, WGPUTextureView uTexMsk, WGPUTextureView uTexDst)
 {
     release();
     const WGPUBindGroupEntry bindGroupEntries[] {
@@ -490,7 +569,51 @@ void WgBindGroupTexComposeBlend::initialize(WGPUDevice device, WGPUQueue queue,
 }
 
 
-void WgBindGroupTexComposeBlend::release()
+void WgBindGroupTexBlend::release()
+{
+    releaseBindGroup(mBindGroup);
+}
+
+///////////////////////////////////////////////////////////////////////////////
+// WgBindGroupTexBlendMask
+///////////////////////////////////////////////////////////////////////////////
+
+WGPUBindGroupLayout WgBindGroupTexBlendMask::getLayout(WGPUDevice device)
+{
+    if (layout) return layout;
+    const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] {
+        makeBindGroupLayoutEntryStorage(0, WGPUStorageTextureAccess_ReadOnly, WGPUTextureFormat_RGBA8Unorm),
+        makeBindGroupLayoutEntryStorage(1, WGPUStorageTextureAccess_ReadOnly, WGPUTextureFormat_RGBA8Unorm),
+        makeBindGroupLayoutEntryStorage(2, WGPUStorageTextureAccess_ReadOnly, WGPUTextureFormat_RGBA8Unorm),
+        makeBindGroupLayoutEntryStorage(3, WGPUStorageTextureAccess_WriteOnly, WGPUTextureFormat_RGBA8Unorm)
+    };
+    layout = createBindGroupLayout(device, bindGroupLayoutEntries, 4);
+    assert(layout);
+    return layout;
+}
+
+
+void WgBindGroupTexBlendMask::releaseLayout()
+{
+    releaseBindGroupLayout(layout);
+}
+
+
+void WgBindGroupTexBlendMask::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexSrc, WGPUTextureView uTexMsk, WGPUTextureView uTexDst, WGPUTextureView uTexTrg)
+{
+    release();
+    const WGPUBindGroupEntry bindGroupEntries[] {
+        makeBindGroupEntryTextureView(0, uTexSrc),
+        makeBindGroupEntryTextureView(1, uTexMsk),
+        makeBindGroupEntryTextureView(2, uTexDst),
+        makeBindGroupEntryTextureView(3, uTexTrg)
+    };
+    mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 4);
+    assert(mBindGroup);
+}
+
+
+void WgBindGroupTexBlendMask::release()
 {
     releaseBindGroup(mBindGroup);
 }
@@ -504,9 +627,10 @@ WGPUBindGroupLayout WgBindGroupTexMaskCompose::getLayout(WGPUDevice device)
     if (layout) return layout;
     const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] {
         makeBindGroupLayoutEntryStorage(0, WGPUStorageTextureAccess_ReadOnly, WGPUTextureFormat_RGBA8Unorm),
-        makeBindGroupLayoutEntryStorage(1, WGPUStorageTextureAccess_ReadWrite, WGPUTextureFormat_RGBA8Unorm)
+        makeBindGroupLayoutEntryStorage(1, WGPUStorageTextureAccess_ReadOnly, WGPUTextureFormat_RGBA8Unorm),
+        makeBindGroupLayoutEntryStorage(2, WGPUStorageTextureAccess_WriteOnly, WGPUTextureFormat_RGBA8Unorm)
     };
-    layout = createBindGroupLayout(device, bindGroupLayoutEntries, 2);
+    layout = createBindGroupLayout(device, bindGroupLayoutEntries, 3);
     assert(layout);
     return layout;
 }
@@ -518,14 +642,15 @@ void WgBindGroupTexMaskCompose::releaseLayout()
 }
 
 
-void WgBindGroupTexMaskCompose::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexMsk0, WGPUTextureView uTexMsk1)
+void WgBindGroupTexMaskCompose::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexMsk0, WGPUTextureView uTexMsk1, WGPUTextureView uTexTrg)
 {
     release();
     const WGPUBindGroupEntry bindGroupEntries[] {
         makeBindGroupEntryTextureView(0, uTexMsk0),
         makeBindGroupEntryTextureView(1, uTexMsk1),
+        makeBindGroupEntryTextureView(2, uTexTrg)
     };
-    mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 2);
+    mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 3);
     assert(mBindGroup);
 }
 
@@ -535,6 +660,50 @@ void WgBindGroupTexMaskCompose::release()
     releaseBindGroup(mBindGroup);
 }
 
+///////////////////////////////////////////////////////////////////////////////
+// WgBindGroupTexCompose
+///////////////////////////////////////////////////////////////////////////////
+
+WGPUBindGroupLayout WgBindGroupTexCompose::getLayout(WGPUDevice device)
+{
+    if (layout) return layout;
+    const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] {
+        makeBindGroupLayoutEntryStorage(0, WGPUStorageTextureAccess_ReadOnly, WGPUTextureFormat_RGBA8Unorm),
+        makeBindGroupLayoutEntryStorage(1, WGPUStorageTextureAccess_ReadOnly, WGPUTextureFormat_RGBA8Unorm),
+        makeBindGroupLayoutEntryStorage(2, WGPUStorageTextureAccess_ReadOnly, WGPUTextureFormat_RGBA8Unorm),
+        makeBindGroupLayoutEntryStorage(3, WGPUStorageTextureAccess_WriteOnly, WGPUTextureFormat_RGBA8Unorm)
+    };
+    layout = createBindGroupLayout(device, bindGroupLayoutEntries, 4);
+    assert(layout);
+    return layout;
+}
+
+
+void WgBindGroupTexCompose::releaseLayout()
+{
+    releaseBindGroupLayout(layout);
+}
+
+
+void WgBindGroupTexCompose::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexSrc, WGPUTextureView uTexMsk, WGPUTextureView uTexDst, WGPUTextureView uTexTrg)
+{
+    release();
+    const WGPUBindGroupEntry bindGroupEntries[] {
+        makeBindGroupEntryTextureView(0, uTexSrc),
+        makeBindGroupEntryTextureView(1, uTexMsk),
+        makeBindGroupEntryTextureView(2, uTexDst),
+        makeBindGroupEntryTextureView(3, uTexTrg)
+    };
+    mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 4);
+    assert(mBindGroup);
+}
+
+
+void WgBindGroupTexCompose::release()
+{
+    releaseBindGroup(mBindGroup);
+}
+
 ///////////////////////////////////////////////////////////////////////////////
 // WgBindGroupOpacity
 ///////////////////////////////////////////////////////////////////////////////

src/renderer/wg_engine/tvgWgBindGroups.h

@@ -119,7 +119,7 @@ struct WgBindGroupTexture : public WgBindGroup
 };
 
 // @group(0 or 1)
-struct WgBindGroupTextureStorageRgba : public WgBindGroup
+struct WgBindGroupTextureStorageRgbaWO : public WgBindGroup
 {
     static WGPUBindGroupLayout layout;
     static WGPUBindGroupLayout getLayout(WGPUDevice device);
@@ -131,7 +131,7 @@ struct WgBindGroupTextureStorageRgba : public WgBindGroup
 };
 
 // @group(0 or 1)
-struct WgBindGroupTextureStorageBgra : public WgBindGroup
+struct WgBindGroupTextureStorageRgbaRO : public WgBindGroup
 {
     static WGPUBindGroupLayout layout;
     static WGPUBindGroupLayout getLayout(WGPUDevice device);
@@ -142,6 +142,37 @@ struct WgBindGroupTextureStorageBgra : public WgBindGroup
     void release();
 };
 
+
+// @group(0 or 1)
+struct WgBindGroupTextureStorageBgraRO : public WgBindGroup
+{
+    static WGPUBindGroupLayout layout;
+    static WGPUBindGroupLayout getLayout(WGPUDevice device);
+    static WGPUBindGroupLayout getLayoutRO(WGPUDevice device);
+    static WGPUBindGroupLayout getLayoutWO(WGPUDevice device);
+    static void releaseLayout();
+
+    void initialize(WGPUDevice device, WGPUQueue queue,
+                    WGPUTextureView uTexture);
+    void release();
+};
+
+
+// @group(0 or 1)
+struct WgBindGroupTextureStorageBgraWO : public WgBindGroup
+{
+    static WGPUBindGroupLayout layout;
+    static WGPUBindGroupLayout getLayout(WGPUDevice device);
+    static WGPUBindGroupLayout getLayoutRO(WGPUDevice device);
+    static WGPUBindGroupLayout getLayoutWO(WGPUDevice device);
+    static void releaseLayout();
+
+    void initialize(WGPUDevice device, WGPUQueue queue,
+                    WGPUTextureView uTexture);
+    void release();
+};
+
+
 // @group(0 or 1)
 struct WgBindGroupTextureSampled : public WgBindGroup
 {
@@ -156,7 +187,22 @@ struct WgBindGroupTextureSampled : public WgBindGroup
 };
 
 // @group(0)
-struct WgBindGroupTexComposeBlend : public WgBindGroup
+struct WgBindGroupTexBlend : public WgBindGroup
+{
+    static WGPUBindGroupLayout layout;
+    static WGPUBindGroupLayout getLayout(WGPUDevice device);
+    static void releaseLayout();
+
+    void initialize(WGPUDevice device, WGPUQueue queue,
+                    WGPUTextureView uTexSrc,
+                    WGPUTextureView uTexDst,
+                    WGPUTextureView uTexTrg);
+    void release();
+};
+
+
+// @group(0)
+struct WgBindGroupTexBlendMask : public WgBindGroup
 {
     static WGPUBindGroupLayout layout;
     static WGPUBindGroupLayout getLayout(WGPUDevice device);
@@ -165,7 +211,8 @@ struct WgBindGroupTexComposeBlend : public WgBindGroup
     void initialize(WGPUDevice device, WGPUQueue queue,
                     WGPUTextureView uTexSrc,
                     WGPUTextureView uTexMsk,
-                    WGPUTextureView uTexDst);
+                    WGPUTextureView uTexDst,
+                    WGPUTextureView uTexTrg);
     void release();
 };
 
@@ -179,7 +226,24 @@ struct WgBindGroupTexMaskCompose : public WgBindGroup
 
     void initialize(WGPUDevice device, WGPUQueue queue,
                     WGPUTextureView uTexMsk0,
-                    WGPUTextureView uTexMsk1);
+                    WGPUTextureView uTexMsk1,
+                    WGPUTextureView uTexTrg);
+    void release();
+};
+
+
+// @group(0)
+struct WgBindGroupTexCompose : public WgBindGroup
+{
+    static WGPUBindGroupLayout layout;
+    static WGPUBindGroupLayout getLayout(WGPUDevice device);
+    static void releaseLayout();
+
+    void initialize(WGPUDevice device, WGPUQueue queue,
+                    WGPUTextureView uTexSrc,
+                    WGPUTextureView uTexMsk,
+                    WGPUTextureView uTexDst,
+                    WGPUTextureView uTexTrg);
     void release();
 };
 

src/renderer/wg_engine/tvgWgCommon.cpp

@@ -21,6 +21,9 @@
  */
 
 #include "tvgWgCommon.h"
+#ifdef __EMSCRIPTEN__
+#include <emscripten.h>
+#endif
 #include <iostream>
 
 //*****************************************************************************
@@ -50,12 +53,13 @@ void WgContext::initialize(WGPUInstance instance, WGPUSurface surface)
     };
     // request adapter
     wgpuInstanceRequestAdapter(instance, &requestAdapterOptions, onAdapterRequestEnded, &adapter);
+    #ifdef __EMSCRIPTEN__
+    while (!adapter) emscripten_sleep(10);
+    #endif
     assert(adapter);
 
-    // adapter enumerate features
+    // get feature names
     size_t featuresCount = wgpuAdapterEnumerateFeatures(adapter, featureNames);
-    wgpuAdapterGetProperties(adapter, &adapterProperties);
-    wgpuAdapterGetLimits(adapter, &supportedLimits);
 
     // request device
     WGPUDeviceDescriptor deviceDesc{};
@@ -76,6 +80,9 @@ void WgContext::initialize(WGPUInstance instance, WGPUSurface surface)
     };
     // request device
     wgpuAdapterRequestDevice(adapter, &deviceDesc, onDeviceRequestEnded, &device);
+    #ifdef __EMSCRIPTEN__
+    while (!device) emscripten_sleep(10);
+    #endif
     assert(device);
 
     // on device error function
@@ -535,8 +542,10 @@ WGPUShaderModule WgPipeline::createShaderModule(WGPUDevice device, const char* c
     WGPUShaderModuleDescriptor shaderModuleDesc{};
     shaderModuleDesc.nextInChain = &shaderModuleWGSLDesc.chain;
     shaderModuleDesc.label = label;
+    #ifndef __EMSCRIPTEN__
     shaderModuleDesc.hintCount = 0;
     shaderModuleDesc.hints = nullptr;
+    #endif
     return wgpuDeviceCreateShaderModule(device, &shaderModuleDesc);
 }
 

src/renderer/wg_engine/tvgWgPipelines.cpp

@@ -298,17 +298,18 @@ void WgPipelineImage::initialize(WGPUDevice device, WgPipelineBlendType blendTyp
 // compute pipelines
 //************************************************************************
 
-void WgPipelineClear::initialize(WGPUDevice device)
+void WgPipelineCopy::initialize(WGPUDevice device)
 {
     // bind groups and layouts
     WGPUBindGroupLayout bindGroupLayouts[] = {
-        WgBindGroupTextureStorageRgba::getLayout(device)
+        WgBindGroupTextureStorageRgbaRO::getLayout(device),
+        WgBindGroupTextureStorageRgbaWO::getLayout(device)
     };
 
     // shader source and labels
-    auto shaderSource = cShaderSource_PipelineComputeClear;
+    auto shaderSource = cShaderSource_PipelineComputeCopy;
-    auto shaderLabel = "The compute shader clear";
+    auto shaderLabel = "The compute shader copy";
-    auto pipelineLabel = "The compute pipeline clear";
+    auto pipelineLabel = "The compute pipeline copy";
 
     // allocate all pipeline handles
     allocate(device,
@@ -321,8 +322,7 @@ void WgPipelineBlend::initialize(WGPUDevice device, const char *shaderSource)
 {
     // bind groups and layouts
     WGPUBindGroupLayout bindGroupLayouts[] = {
-        WgBindGroupTextureStorageRgba::getLayout(device),
+        WgBindGroupTexBlend::getLayout(device),
-        WgBindGroupTextureStorageRgba::getLayout(device),
         WgBindGroupBlendMethod::getLayout(device),
         WgBindGroupOpacity::getLayout(device)
     };
@@ -342,7 +342,7 @@ void WgPipelineBlendMask::initialize(WGPUDevice device, const char *shaderSource
 {
     // bind groups and layouts
     WGPUBindGroupLayout bindGroupLayouts[] = {
-        WgBindGroupTexComposeBlend::getLayout(device),
+        WgBindGroupTexBlendMask::getLayout(device),
         WgBindGroupBlendMethod::getLayout(device),
         WgBindGroupOpacity::getLayout(device)
     };
@@ -381,7 +381,7 @@ void WgPipelineCompose::initialize(WGPUDevice device)
 {
     // bind groups and layouts
     WGPUBindGroupLayout bindGroupLayouts[] = {
-        WgBindGroupTexComposeBlend::getLayout(device),
+        WgBindGroupTexCompose::getLayout(device),
         WgBindGroupCompositeMethod::getLayout(device),
         WgBindGroupBlendMethod::getLayout(device),
         WgBindGroupOpacity::getLayout(device)
@@ -403,8 +403,8 @@ void WgPipelineAntiAliasing::initialize(WGPUDevice device)
 {
     // bind groups and layouts
     WGPUBindGroupLayout bindGroupLayouts[] = {
-        WgBindGroupTextureStorageRgba::getLayout(device),
+        WgBindGroupTextureStorageRgbaRO::getLayout(device),
-        WgBindGroupTextureStorageBgra::getLayout(device)
+        WgBindGroupTextureStorageBgraWO::getLayout(device)
     };
 
     // shader source and labels
@@ -436,7 +436,7 @@ void WgPipelines::initialize(WgContext& context)
         image[type].initialize(context.device, (WgPipelineBlendType)type);
     }
     // compute pipelines
-    computeClear.initialize(context.device);
+    computeCopy.initialize(context.device);
     computeBlendSolid.initialize(context.device, cShaderSource_PipelineComputeBlendSolid);
     computeBlendGradient.initialize(context.device, cShaderSource_PipelineComputeBlendGradient);
     computeBlendImage.initialize(context.device, cShaderSource_PipelineComputeBlendImage);
@@ -453,11 +453,14 @@ void WgPipelines::initialize(WgContext& context)
 
 void WgPipelines::release()
 {
+    WgBindGroupTexCompose::releaseLayout();
     WgBindGroupTexMaskCompose::releaseLayout();
-    WgBindGroupTexComposeBlend::releaseLayout();
+    WgBindGroupTexBlendMask::releaseLayout();
-    WgBindGroupTextureSampled::releaseLayout();
+    WgBindGroupTexBlend::releaseLayout();
-    WgBindGroupTextureStorageBgra::releaseLayout();
+    WgBindGroupTextureStorageBgraRO::releaseLayout();
-    WgBindGroupTextureStorageRgba::releaseLayout();
+    WgBindGroupTextureStorageBgraWO::releaseLayout();
+    WgBindGroupTextureStorageRgbaRO::releaseLayout();
+    WgBindGroupTextureStorageRgbaWO::releaseLayout();
     WgBindGroupTexture::releaseLayout();
     WgBindGroupOpacity::releaseLayout();
     WgBindGroupPicture::releaseLayout();
@@ -469,14 +472,13 @@ void WgPipelines::release()
     // compute pipelines
     computeAntiAliasing.release();
     computeCompose.release();
-    computeMaskCompose.release();
     computeBlendImageMask.release();
     computeBlendGradientMask.release();
     computeBlendSolidMask.release();
     computeBlendImage.release();
     computeBlendGradient.release();
     computeBlendSolid.release();
-    computeClear.release();
+    computeCopy.release();
     // fill pipelines
     for (uint8_t type = (uint8_t)WgPipelineBlendType::SrcOver; type <= (uint8_t)WgPipelineBlendType::Custom; type++) {
         image[type].release();

src/renderer/wg_engine/tvgWgPipelines.h

@@ -129,13 +129,14 @@ struct WgPipelineImage: public WgRenderPipeline
 // compute pipelines
 //*****************************************************************************
 
-struct WgPipelineClear: public WgComputePipeline
+struct WgPipelineCopy: public WgComputePipeline
 {
     void initialize(WGPUDevice device) override;
-    void use(WGPUComputePassEncoder encoder, WgBindGroupTextureStorageRgba& groupTexDst)
+    void use(WGPUComputePassEncoder encoder, WgBindGroupTextureStorageRgbaRO& groupTexSrc, WgBindGroupTextureStorageRgbaWO& groupTexDst)
     {
         set(encoder);
-        groupTexDst.set(encoder, 0);
+        groupTexSrc.set(encoder, 0);
+        groupTexDst.set(encoder, 1);
     }
 };
 
@@ -144,13 +145,12 @@ struct WgPipelineBlend: public WgComputePipeline
 {
     void initialize(WGPUDevice device) override { assert(false); };
     void initialize(WGPUDevice device, const char *shaderSource);
-    void use(WGPUComputePassEncoder encoder, WgBindGroupTextureStorageRgba& groupTexSrc, WgBindGroupTextureStorageRgba& groupTexDst, WgBindGroupBlendMethod& blendMethod, WgBindGroupOpacity& groupOpacity)
+    void use(WGPUComputePassEncoder encoder, WgBindGroupTexBlend& groupTexBlend, WgBindGroupBlendMethod& blendMethod, WgBindGroupOpacity& groupOpacity)
     {
         set(encoder);
-        groupTexSrc.set(encoder, 0);
+        groupTexBlend.set(encoder, 0);
-        groupTexDst.set(encoder, 1);
+        blendMethod.set(encoder, 1);
-        blendMethod.set(encoder, 2);
+        groupOpacity.set(encoder, 2);
-        groupOpacity.set(encoder, 3);
     }
 };
 
@@ -159,10 +159,10 @@ 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)
+    void use(WGPUComputePassEncoder encoder, WgBindGroupTexBlendMask& groupTexBlendMask, WgBindGroupBlendMethod& blendMethod, WgBindGroupOpacity& groupOpacity)
     {
         set(encoder);
-        groupTexs.set(encoder, 0);
+        groupTexBlendMask.set(encoder, 0);
         blendMethod.set(encoder, 1);
         groupOpacity.set(encoder, 2);
     }
@@ -172,10 +172,10 @@ struct WgPipelineBlendMask: public WgComputePipeline
 struct WgPipelineMaskCompose: public WgComputePipeline
 {
     void initialize(WGPUDevice device) override;
-    void use(WGPUComputePassEncoder encoder, WgBindGroupTexMaskCompose& groupTexs)
+    void use(WGPUComputePassEncoder encoder, WgBindGroupTexMaskCompose& groupTexMaskCompose)
     {
         set(encoder);
-        groupTexs.set(encoder, 0);
+        groupTexMaskCompose.set(encoder, 0);
     }
 };
 
@@ -183,10 +183,10 @@ struct WgPipelineMaskCompose: public WgComputePipeline
 struct WgPipelineCompose: public WgComputePipeline
 {
     void initialize(WGPUDevice device) override;
-    void use(WGPUComputePassEncoder encoder, WgBindGroupTexComposeBlend& groupTexs, WgBindGroupCompositeMethod& groupComposeMethod, WgBindGroupBlendMethod& groupBlendMethod, WgBindGroupOpacity& groupOpacity)
+    void use(WGPUComputePassEncoder encoder, WgBindGroupTexCompose& groupTexCompose, WgBindGroupCompositeMethod& groupComposeMethod, WgBindGroupBlendMethod& groupBlendMethod, WgBindGroupOpacity& groupOpacity)
     {
         set(encoder);
-        groupTexs.set(encoder, 0);
+        groupTexCompose.set(encoder, 0);
         groupComposeMethod.set(encoder, 1);
         groupBlendMethod.set(encoder, 2);
         groupOpacity.set(encoder, 3);
@@ -197,7 +197,7 @@ struct WgPipelineCompose: public WgComputePipeline
 struct WgPipelineAntiAliasing: public WgComputePipeline
 {
     void initialize(WGPUDevice device) override;
-    void use(WGPUComputePassEncoder encoder, WgBindGroupTextureStorageRgba& groupTexSrc, WgBindGroupTextureStorageBgra& groupTexDst)
+    void use(WGPUComputePassEncoder encoder, WgBindGroupTextureStorageRgbaRO& groupTexSrc, WgBindGroupTextureStorageBgraWO& groupTexDst)
     {
         set(encoder);
         groupTexSrc.set(encoder, 0);
@@ -222,7 +222,7 @@ struct WgPipelines
     WgPipelineRadial radial[3];
     WgPipelineImage image[3];
     // compute pipelines
-    WgPipelineClear computeClear;
+    WgPipelineCopy computeCopy;
     WgPipelineBlend computeBlendSolid;
     WgPipelineBlend computeBlendGradient;
     WgPipelineBlend computeBlendImage;

src/renderer/wg_engine/tvgWgRenderTarget.cpp

@@ -54,10 +54,12 @@
     assert(texViewColor);
     assert(texViewStencil);
     // initialize bind group for blitting
-    if (format == WGPUTextureFormat_RGBA8Unorm)
+    if (format == WGPUTextureFormat_RGBA8Unorm) {
-        bindGroupTexStorageRgba.initialize(context.device, context.queue, texViewColor);
+        bindGroupTexStorageRgbaRO.initialize(context.device, context.queue, texViewColor);
+        bindGroupTexStorageRgbaWO.initialize(context.device, context.queue, texViewColor);
+    }
     if (format == WGPUTextureFormat_BGRA8Unorm)
-        bindGroupTexStorageBgra.initialize(context.device, context.queue, texViewColor);
+        bindGroupTexStorageBgraWO.initialize(context.device, context.queue, texViewColor);
     // initialize window binding groups
     WgShaderTypeMat4x4f viewMat(w, h);
     mBindGroupCanvas.initialize(context.device, context.queue, viewMat);
@@ -69,7 +71,9 @@ void WgRenderStorage::release(WgContext& context)
 {
     mRenderPassEncoder = nullptr;
     mBindGroupCanvas.release();
-    bindGroupTexStorageRgba.release();
+    bindGroupTexStorageBgraWO.release();
+    bindGroupTexStorageRgbaWO.release();
+    bindGroupTexStorageRgbaRO.release();
     context.releaseTextureView(texViewStencil);
     context.releaseTextureView(texViewColor);
     context.releaseTexture(texStencil);
@@ -210,29 +214,24 @@ void WgRenderStorage::drawPictureClipPath(WgContext& context, WgRenderDataPictur
 }
 
 
-void WgRenderStorage::clear(WGPUCommandEncoder commandEncoder)
-{
-    assert(commandEncoder);
-    WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder);
-    mPipelines->computeClear.use(computePassEncoder, bindGroupTexStorageRgba);
-    dispatchWorkgroups(computePassEncoder);
-    endComputePass(computePassEncoder);
-}
-
 
 void WgRenderStorage::blend(
     WgContext& context,
     WGPUCommandEncoder commandEncoder,
     WgPipelineBlend* pipeline,
-    WgRenderStorage* targetSrc,
+    WgRenderStorage* texSrc,
+    WgRenderStorage* texDst,
     WgBindGroupBlendMethod* blendMethod,
     WgBindGroupOpacity* opacity)
 {
     assert(commandEncoder);
-    assert(targetSrc);
+    assert(texSrc);
+    assert(texDst);
     assert(pipeline);
+    WgBindGroupTexBlend texBlend;
+    texBlend.initialize(context.device, context.queue, texSrc->texViewColor, texDst->texViewColor, texViewColor);
     WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder);
-    pipeline->use(computePassEncoder, targetSrc->bindGroupTexStorageRgba, bindGroupTexStorageRgba, *blendMethod, *opacity);
+    pipeline->use(computePassEncoder, texBlend, *blendMethod, *opacity);
     dispatchWorkgroups(computePassEncoder);
     endComputePass(computePassEncoder);
 }
@@ -244,31 +243,33 @@ void WgRenderStorage::blendMask(
     WgPipelineBlendMask* pipeline,
     WgRenderStorage* texMsk,
     WgRenderStorage* texSrc,
+    WgRenderStorage* texDst,
     WgBindGroupBlendMethod* blendMethod,
     WgBindGroupOpacity* opacity)
 {
     assert(commandEncoder);
     assert(texSrc);
     assert(texMsk);
-    WgBindGroupTexComposeBlend composeBlend;
+    WgBindGroupTexBlendMask texBlendMask;
-    composeBlend.initialize(context.device, context.queue, texSrc->texViewColor, texMsk->texViewColor, texViewColor);
+    texBlendMask.initialize(context.device, context.queue, texSrc->texViewColor, texMsk->texViewColor, texDst->texViewColor, texViewColor);
     WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder);
-    pipeline->use(computePassEncoder, composeBlend, *blendMethod, *opacity);
+    pipeline->use(computePassEncoder, texBlendMask, *blendMethod, *opacity);
     dispatchWorkgroups(computePassEncoder);
     endComputePass(computePassEncoder);
-    composeBlend.release();
+    texBlendMask.release();
 };
 
 
 void WgRenderStorage::maskCompose(
         WgContext& context,
         WGPUCommandEncoder commandEncoder,
-        WgRenderStorage* texMsk0)
+        WgRenderStorage* texMsk0,
+        WgRenderStorage* texMsk1)
 {
     assert(commandEncoder);
     assert(texMsk0);
     WgBindGroupTexMaskCompose maskCompose;
-    maskCompose.initialize(context.device, context.queue, texMsk0->texViewColor, texViewColor);
+    maskCompose.initialize(context.device, context.queue, texMsk0->texViewColor, texMsk1->texViewColor, texViewColor);
     WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder);
     mPipelines->computeMaskCompose.use(computePassEncoder, maskCompose);
     dispatchWorkgroups(computePassEncoder);
@@ -282,6 +283,7 @@ void WgRenderStorage::compose(
     WGPUCommandEncoder commandEncoder,
     WgRenderStorage* texSrc,
     WgRenderStorage* texMsk,
+    WgRenderStorage* texDst,
     WgBindGroupCompositeMethod* composeMethod,
     WgBindGroupBlendMethod* blendMethod,
     WgBindGroupOpacity* opacity)
@@ -289,13 +291,13 @@ void WgRenderStorage::compose(
     assert(commandEncoder);
     assert(texSrc);
     assert(texMsk);
-    WgBindGroupTexComposeBlend composeBlend;
+    WgBindGroupTexCompose texCompose;
-    composeBlend.initialize(context.device, context.queue, texSrc->texViewColor, texMsk->texViewColor, texViewColor);
+    texCompose.initialize(context.device, context.queue, texSrc->texViewColor, texMsk->texViewColor, texDst->texViewColor, texViewColor);
     WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder);
-    mPipelines->computeCompose.use(computePassEncoder, composeBlend, *composeMethod, *blendMethod, *opacity);
+    mPipelines->computeCompose.use(computePassEncoder, texCompose, *composeMethod, *blendMethod, *opacity);
     dispatchWorkgroups(computePassEncoder);
     endComputePass(computePassEncoder);
-    composeBlend.release();
+    texCompose.release();
 }
 
 
@@ -304,11 +306,20 @@ void WgRenderStorage::antialias(WGPUCommandEncoder commandEncoder, WgRenderStora
     assert(commandEncoder);
     assert(targetSrc);
     WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder);
-    mPipelines->computeAntiAliasing.use(computePassEncoder, targetSrc->bindGroupTexStorageRgba, bindGroupTexStorageBgra);
+    mPipelines->computeAntiAliasing.use(computePassEncoder, targetSrc->bindGroupTexStorageRgbaRO, bindGroupTexStorageBgraWO);
     dispatchWorkgroups(computePassEncoder);
     endComputePass(computePassEncoder);
 }
 
+void WgRenderStorage::copy(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc)
+{
+    assert(commandEncoder);
+    assert(targetSrc);
+    WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder);
+    mPipelines->computeCopy.use(computePassEncoder, targetSrc->bindGroupTexStorageRgbaRO, bindGroupTexStorageRgbaWO);
+    dispatchWorkgroups(computePassEncoder);
+    endComputePass(computePassEncoder);
+}
 
 void WgRenderStorage::dispatchWorkgroups(WGPUComputePassEncoder computePassEncoder)
 {
@@ -323,8 +334,8 @@ void WgRenderStorage::beginRenderPass(WGPUCommandEncoder commandEncoder, bool cl
     // render pass depth stencil attachment
     WGPURenderPassDepthStencilAttachment depthStencilAttachment{};
     depthStencilAttachment.view = texViewStencil;
-    depthStencilAttachment.depthLoadOp = WGPULoadOp_Load;
+    depthStencilAttachment.depthLoadOp = WGPULoadOp_Undefined;
-    depthStencilAttachment.depthStoreOp = WGPUStoreOp_Discard;
+    depthStencilAttachment.depthStoreOp = WGPUStoreOp_Undefined;
     depthStencilAttachment.depthClearValue = 1.0f;
     depthStencilAttachment.depthReadOnly = false;
     depthStencilAttachment.stencilLoadOp = WGPULoadOp_Clear;
@@ -334,6 +345,9 @@ void WgRenderStorage::beginRenderPass(WGPUCommandEncoder commandEncoder, bool cl
     // render pass color attachment
     WGPURenderPassColorAttachment colorAttachment{};
     colorAttachment.view = texViewColor;
+    #ifdef __EMSCRIPTEN__
+    colorAttachment.depthSlice = WGPU_DEPTH_SLICE_UNDEFINED;
+    #endif
     colorAttachment.resolveTarget = nullptr;
     colorAttachment.loadOp = clear ? WGPULoadOp_Clear : WGPULoadOp_Load;
     colorAttachment.clearValue = { 0, 0, 0, 0 };

src/renderer/wg_engine/tvgWgRenderTarget.h

@@ -36,8 +36,9 @@ public:
     WGPUTexture texStencil{};
     WGPUTextureView texViewColor{};
     WGPUTextureView texViewStencil{};
-    WgBindGroupTextureStorageRgba bindGroupTexStorageRgba;
+    WgBindGroupTextureStorageRgbaRO bindGroupTexStorageRgbaRO;
-    WgBindGroupTextureStorageBgra bindGroupTexStorageBgra;
+    WgBindGroupTextureStorageRgbaWO bindGroupTexStorageRgbaWO;
+    WgBindGroupTextureStorageBgraWO bindGroupTexStorageBgraWO;
     uint32_t samples{};
     uint32_t width{};
     uint32_t height{};
@@ -54,12 +55,12 @@ public:
     void renderPicture(WgContext& context, WgRenderDataPicture* renderData, WgPipelineBlendType blendType);
     void renderClipPath(WgContext& context, WgRenderDataPaint* renderData);
 
-    void clear(WGPUCommandEncoder commandEncoder);
     void blend(
         WgContext& context,
         WGPUCommandEncoder commandEncoder,
         WgPipelineBlend* pipeline,
         WgRenderStorage* targetSrc,
+        WgRenderStorage* targetDst,
         WgBindGroupBlendMethod* blendMethod,
         WgBindGroupOpacity* opacity);
     void blendMask(
@@ -68,21 +69,25 @@ public:
         WgPipelineBlendMask* pipeline,
         WgRenderStorage* texMsk,
         WgRenderStorage* texSrc,
+        WgRenderStorage* texDst,
         WgBindGroupBlendMethod* blendMethod,
         WgBindGroupOpacity* opacity);
     void maskCompose(
         WgContext& context,
         WGPUCommandEncoder commandEncoder,
-        WgRenderStorage* texMsk0);
+        WgRenderStorage* texMsk0,
+        WgRenderStorage* texMsk1);
     void compose(
         WgContext& context,
         WGPUCommandEncoder commandEncoder,
         WgRenderStorage* texMsk,
         WgRenderStorage* texSrc,
+        WgRenderStorage* texDst,
         WgBindGroupCompositeMethod* composeMethod,
         WgBindGroupBlendMethod* blendMethod,
         WgBindGroupOpacity* opacity);
     void antialias(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc);
+    void copy(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc);
 private:
     void drawShape(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType);
     void drawStroke(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType);

src/renderer/wg_engine/tvgWgRenderer.cpp

@@ -68,6 +68,7 @@ void WgRenderer::release()
     mOpacityPool.release(mContext);
     mRenderStorageRoot.release(mContext);
     mRenderStorageMask.release(mContext);
+    mRenderStorageCopy.release(mContext);
     mRenderStorageScreen.release(mContext);
     mRenderStorageInterm.release(mContext);
     mPipelines.release();
@@ -190,7 +191,8 @@ void WgRenderer::renderClipPath(Array<WgRenderDataPaint*>& clips)
         mRenderStorageInterm.beginRenderPass(mCommandEncoder, true);
         mRenderStorageInterm.renderClipPath(mContext, clips[i]);
         mRenderStorageInterm.endRenderPass();
-        mRenderStorageMask.maskCompose(mContext, mCommandEncoder, &mRenderStorageInterm);
+        mRenderStorageCopy.copy(mCommandEncoder, &mRenderStorageMask);
+        mRenderStorageMask.maskCompose(mContext, mCommandEncoder, &mRenderStorageInterm, &mRenderStorageCopy);
     }
 }
 
@@ -222,8 +224,9 @@ bool WgRenderer::renderShape(RenderData data)
         WgPipelineBlendMask* pipeline = &mContext.pipelines->computeBlendSolidMask;
         if (dataShape->renderSettingsShape.fillType != WgRenderSettingsType::Solid)
             pipeline = &mContext.pipelines->computeBlendGradientMask;
+        mRenderStorageCopy.copy(mCommandEncoder, renderStorage);
         renderStorage->blendMask(mContext, mCommandEncoder, 
-            pipeline, &mRenderStorageMask, &mRenderStorageInterm, blendMethod, opacity);
+            pipeline, &mRenderStorageMask, &mRenderStorageInterm, &mRenderStorageCopy, blendMethod, opacity);
         // restore current render pass
         renderStorage->beginRenderPass(mCommandEncoder, false);
     // use hardware blend
@@ -243,8 +246,9 @@ bool WgRenderer::renderShape(RenderData data)
         WgPipelineBlend* pipeline = &mContext.pipelines->computeBlendSolid;
         if (dataShape->renderSettingsShape.fillType != WgRenderSettingsType::Solid)
             pipeline = &mContext.pipelines->computeBlendGradient;
+        mRenderStorageCopy.copy(mCommandEncoder, renderStorage);
         renderStorage->blend(mContext, mCommandEncoder, 
-            pipeline, &mRenderStorageInterm, blendMethod, opacity);
+            pipeline, &mRenderStorageInterm, &mRenderStorageCopy, blendMethod, opacity);
         // restore current render pass
         renderStorage->beginRenderPass(mCommandEncoder, false);
     }
@@ -277,8 +281,9 @@ bool WgRenderer::renderImage(RenderData data)
         WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod);
         WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, 255);
         WgPipelineBlendMask* pipeline = &mContext.pipelines->computeBlendImageMask;
+        mRenderStorageCopy.copy(mCommandEncoder, renderStorage);
         renderStorage->blendMask(mContext, mCommandEncoder, 
-            pipeline, &mRenderStorageMask, &mRenderStorageInterm, blendMethod, opacity);
+            pipeline, &mRenderStorageMask, &mRenderStorageInterm, &mRenderStorageCopy, blendMethod, opacity);
         // restore current render pass
         renderStorage->beginRenderPass(mCommandEncoder, false);
     // use hardware blend
@@ -296,8 +301,9 @@ bool WgRenderer::renderImage(RenderData data)
         WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod);
         WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, 255);
         WgPipelineBlend* pipeline = &mContext.pipelines->computeBlendImage;
+        mRenderStorageCopy.copy(mCommandEncoder, renderStorage);
         renderStorage->blend(mContext, mCommandEncoder, 
-            pipeline, &mRenderStorageInterm, blendMethod, opacity);
+            pipeline, &mRenderStorageInterm, &mRenderStorageCopy, blendMethod, opacity);
         // restore current render pass
         renderStorage->beginRenderPass(mCommandEncoder, false);
     }
@@ -418,12 +424,18 @@ bool WgRenderer::target(WGPUInstance instance, WGPUSurface surface, uint32_t w,
     surfaceConfiguration.alphaMode = WGPUCompositeAlphaMode_Auto;
     surfaceConfiguration.width = w;
     surfaceConfiguration.height = h;
+    #ifdef __EMSCRIPTEN__
+    surfaceConfiguration.presentMode = WGPUPresentMode_Fifo;
+    #else
     surfaceConfiguration.presentMode = WGPUPresentMode_Immediate;
+    #endif
     wgpuSurfaceConfigure(mContext.surface, &surfaceConfiguration);
 
     initialize();
+
     mRenderStorageInterm.initialize(mContext, w, h, WG_SSAA_SAMPLES);
     mRenderStorageMask.initialize(mContext, w, h, WG_SSAA_SAMPLES);
+    mRenderStorageCopy.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;
@@ -469,9 +481,10 @@ 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);
+        mRenderStorageCopy.copy(mCommandEncoder, mRenderStorageStack.last());
         mRenderStorageStack.last()->blend(mContext, mCommandEncoder,
             &mContext.pipelines->computeBlendImage,
-            renderStorageSrc, blendMethod, opacity);
+            renderStorageSrc, &mRenderStorageCopy, blendMethod, opacity);
 
         // back render targets to the pool
         mRenderStoragePool.free(mContext, renderStorageSrc);
@@ -495,8 +508,9 @@ bool WgRenderer::endComposite(TVG_UNUSED Compositor* cmp)
 
         // compose and blend
         // dest = blend(dest, compose(src, msk, composeMethod), blendMethod, opacity)
+        mRenderStorageCopy.copy(mCommandEncoder, renderStorageDst);
         renderStorageDst->compose(mContext, mCommandEncoder,
-            renderStorageSrc, renderStorageMsk,
+            renderStorageSrc, renderStorageMsk, &mRenderStorageCopy,
             composeMethod, blendMethod, opacity);
 
         // back render targets to the pool

src/renderer/wg_engine/tvgWgRenderer.h

@@ -66,6 +66,7 @@ private:
 
     WGPUCommandEncoder mCommandEncoder{};
     WgRenderStorage mRenderStorageInterm; // intermidiate buffer to render
+    WgRenderStorage mRenderStorageCopy;   // copy of destination target (blend and compostition)
     WgRenderStorage mRenderStorageMask;   // buffer to render mask 
     WgRenderStorage mRenderStorageRoot;   // root render storage
     WgRenderStorage mRenderStorageScreen; // storage with data after antializing

src/renderer/wg_engine/tvgWgShaderSrc.cpp

@@ -369,10 +369,11 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f {
 //************************************************************************
 
 const std::string strBlendShaderHeader = WG_SHADER_SOURCE(
-@group(0) @binding(0) var imageSrc : texture_storage_2d<rgba8unorm, read_write>;
+@group(0) @binding(0) var imageSrc : texture_storage_2d<rgba8unorm, read>;
-@group(1) @binding(0) var imageDst : texture_storage_2d<rgba8unorm, read_write>;
+@group(0) @binding(1) var imageDst : texture_storage_2d<rgba8unorm, read>;
-@group(2) @binding(0) var<uniform> blendMethod : u32;
+@group(0) @binding(2) var imageTrg : texture_storage_2d<rgba8unorm, write>;
-@group(3) @binding(0) var<uniform> opacity : f32;
+@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) {
@@ -396,7 +397,8 @@ fn cs_main( @builtin(global_invocation_id) id: vec3u) {
 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(0) @binding(2) var imageDst : texture_storage_2d<rgba8unorm, read>;
+@group(0) @binding(3) var imageTrg : texture_storage_2d<rgba8unorm, write>;
 @group(1) @binding(0) var<uniform> blendMethod : u32;
 @group(2) @binding(0) var<uniform> opacity : f32;
 
@@ -483,7 +485,7 @@ const std::string strBlendShaderPostConditionsImage = WG_SHADER_SOURCE(
 );
 
 const std::string strBlendShaderFooter = WG_SHADER_SOURCE(
-    textureStore(imageDst, id.xy, vec4(Rc, Ra));
+    textureStore(imageTrg, id.xy, vec4(Rc, Ra));
 }
 );
 
@@ -539,7 +541,7 @@ const char* cShaderSource_PipelineComputeBlendImageMask = strComputeBlendImageMa
 
 // pipeline shader modules clear
 const char* cShaderSource_PipelineComputeClear = WG_SHADER_SOURCE(
-@group(0) @binding(0) var imageDst : texture_storage_2d<rgba8unorm, read_write>;
+@group(0) @binding(0) var imageDst : texture_storage_2d<rgba8unorm, write>;
 
 @compute @workgroup_size(8, 8)
 fn cs_main( @builtin(global_invocation_id) id: vec3u) {
@@ -547,16 +549,30 @@ fn cs_main( @builtin(global_invocation_id) id: vec3u) {
 }
 );
 
+
+// pipeline shader modules copy
+const char* cShaderSource_PipelineComputeCopy = WG_SHADER_SOURCE(
+@group(0) @binding(0) var imageSrc : texture_storage_2d<rgba8unorm, read>;
+@group(1) @binding(0) var imageDst : texture_storage_2d<rgba8unorm, write>;
+
+@compute @workgroup_size(8, 8)
+fn cs_main( @builtin(global_invocation_id) id: vec3u) {
+    textureStore(imageDst, id.xy, textureLoad(imageSrc, id.xy));
+}
+);
+
+
 // pipeline shader modules compose
 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>;
+@group(0) @binding(1) var imageMsk1 : texture_storage_2d<rgba8unorm, read>;
+@group(0) @binding(2) 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(imageMsk1, id.xy, colorMsk0 * colorMsk1);
+    textureStore(imageTrg, id.xy, colorMsk0 * colorMsk1);
 }
 );
 
@@ -564,7 +580,8 @@ fn cs_main( @builtin(global_invocation_id) id: vec3u) {
 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>;
+@group(0) @binding(2) var imageDst : texture_storage_2d<rgba8unorm, read>;
+@group(0) @binding(3) var imageTrg : texture_storage_2d<rgba8unorm, write>;
 @group(1) @binding(0) var<uniform> composeMethod : u32;
 @group(2) @binding(0) var<uniform> blendMethod : u32;
 @group(3) @binding(0) var<uniform> opacity : f32;
@@ -615,13 +632,13 @@ fn cs_main( @builtin(global_invocation_id) id: vec3u) {
         Ra = Sa;
     }
 
-    textureStore(imageDst, id.xy, vec4f(Rc, Ra));
+    textureStore(imageTrg, id.xy, vec4f(Rc, Ra));
 }
 );
 
 // pipeline shader modules anti-aliasing
 const char* cShaderSource_PipelineComputeAntiAlias = WG_SHADER_SOURCE(
-@group(0) @binding(0) var imageSrc : texture_storage_2d<rgba8unorm, read_write>;
+@group(0) @binding(0) var imageSrc : texture_storage_2d<rgba8unorm, read>;
 @group(1) @binding(0) var imageDst : texture_storage_2d<bgra8unorm, write>;
 
 @compute @workgroup_size(8, 8)

src/renderer/wg_engine/tvgWgShaderSrc.h

@@ -42,6 +42,7 @@ extern const char* cShaderSource_PipelineImage;
 
 // pipeline shader modules clear, compose and blend
 extern const char* cShaderSource_PipelineComputeClear;
+extern const char* cShaderSource_PipelineComputeCopy;
 extern const char* cShaderSource_PipelineComputeBlendSolid;
 extern const char* cShaderSource_PipelineComputeBlendGradient;
 extern const char* cShaderSource_PipelineComputeBlendImage;