wg_engine: scene blending optimization

- used hardware blending stage for scene blending
- used AABB for scene blending
- reduced number of offfscreen buffers coping
- reduced number of render pass switching
- used render pipelines abilities to convert offscreen pixel format to screen pixel format
- removed unused shaders

src/renderer/wg_engine/tvgWgBindGroups.cpp

@@ -45,15 +45,6 @@ WGPUBindGroup WgBindGroupLayouts::createBindGroupTexSampledBuff1Un(WGPUSampler s
 }
 
 
-WGPUBindGroup WgBindGroupLayouts::createBindGroupScreen1WO(WGPUTextureView texView) {
-    const WGPUBindGroupEntry bindGroupEntrys[] = {
-        { .binding = 0, .textureView = texView }
-    };
-    const WGPUBindGroupDescriptor bindGroupDesc { .layout = layoutTexScreen1WO, .entryCount = 1, .entries = bindGroupEntrys };
-    return wgpuDeviceCreateBindGroup(device, &bindGroupDesc);
-}
-
-
 WGPUBindGroup WgBindGroupLayouts::createBindGroupStrorage1WO(WGPUTextureView texView)
 {
     const WGPUBindGroupEntry bindGroupEntrys[] = {
@@ -150,7 +141,6 @@ void WgBindGroupLayouts::initialize(WgContext& context)
     const WGPUTextureBindingLayout texture = { .sampleType = WGPUTextureSampleType_Float, .viewDimension = WGPUTextureViewDimension_2D };
     const WGPUStorageTextureBindingLayout storageTextureWO { .access = WGPUStorageTextureAccess_WriteOnly, .format = WGPUTextureFormat_RGBA8Unorm, .viewDimension = WGPUTextureViewDimension_2D };
     const WGPUStorageTextureBindingLayout storageTextureRO { .access = WGPUStorageTextureAccess_ReadOnly,  .format = WGPUTextureFormat_RGBA8Unorm, .viewDimension = WGPUTextureViewDimension_2D };
-    const WGPUStorageTextureBindingLayout storageScreenWO  { .access = WGPUStorageTextureAccess_WriteOnly, .format = WGPUTextureFormat_BGRA8Unorm, .viewDimension = WGPUTextureViewDimension_2D };
     const WGPUBufferBindingLayout bufferUniform { .type = WGPUBufferBindingType_Uniform };
 
     { // bind group layout tex sampled
@@ -174,15 +164,6 @@ void WgBindGroupLayouts::initialize(WgContext& context)
         assert(layoutTexSampledBuff1Un);
     }
 
-    { // bind group layout tex screen 1 RO
-        const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] {
-            { .binding = 0, .visibility = visibility_frag, .storageTexture = storageScreenWO }
-        };
-        const WGPUBindGroupLayoutDescriptor bindGroupLayoutDesc { .entryCount = 1, .entries = bindGroupLayoutEntries };
-        layoutTexScreen1WO = wgpuDeviceCreateBindGroupLayout(device, &bindGroupLayoutDesc);
-        assert(layoutTexScreen1WO);
-    }
-
     { // bind group layout tex storage 1 WO
         const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] {
             { .binding = 0, .visibility = visibility_frag, .storageTexture = storageTextureWO }
@@ -259,7 +240,6 @@ void WgBindGroupLayouts::release(WgContext& context)
     wgpuBindGroupLayoutRelease(layoutBuffer3Un);
     wgpuBindGroupLayoutRelease(layoutBuffer2Un);
     wgpuBindGroupLayoutRelease(layoutBuffer1Un);
-    wgpuBindGroupLayoutRelease(layoutTexScreen1WO);
     wgpuBindGroupLayoutRelease(layoutTexStrorage3RO);
     wgpuBindGroupLayoutRelease(layoutTexStrorage2RO);
     wgpuBindGroupLayoutRelease(layoutTexStrorage1RO);

src/renderer/wg_engine/tvgWgBindGroups.h

@@ -31,7 +31,6 @@ private:
 public:
     WGPUBindGroupLayout layoutTexSampled{};
     WGPUBindGroupLayout layoutTexSampledBuff1Un{};
-    WGPUBindGroupLayout layoutTexScreen1WO{};
     WGPUBindGroupLayout layoutTexStrorage1WO{};
     WGPUBindGroupLayout layoutTexStrorage1RO{};
     WGPUBindGroupLayout layoutTexStrorage2RO{};
@@ -42,7 +41,6 @@ public:
 public:
     WGPUBindGroup createBindGroupTexSampled(WGPUSampler sampler, WGPUTextureView texView);
     WGPUBindGroup createBindGroupTexSampledBuff1Un(WGPUSampler sampler, WGPUTextureView texView, WGPUBuffer buff);
-    WGPUBindGroup createBindGroupScreen1WO(WGPUTextureView texView);
     WGPUBindGroup createBindGroupStrorage1WO(WGPUTextureView texView);
     WGPUBindGroup createBindGroupStrorage1RO(WGPUTextureView texView);
     WGPUBindGroup createBindGroupStrorage2RO(WGPUTextureView texView0, WGPUTextureView texView1);

src/renderer/wg_engine/tvgWgCommon.cpp

@@ -48,7 +48,7 @@ void WgContext::initialize(WGPUInstance instance, WGPUSurface surface)
     // get adapter and surface properties
     WGPUFeatureName featureNames[32]{};
     size_t featuresCount = wgpuAdapterEnumerateFeatures(adapter, featureNames);
-    preferredFormat = wgpuSurfaceGetPreferredFormat(surface, adapter);
+    preferredFormat = WGPUTextureFormat_BGRA8Unorm;
 
     // request device
     const WGPUDeviceDescriptor deviceDesc { .nextInChain = nullptr, .label = "The device", .requiredFeatureCount = featuresCount, .requiredFeatures = featureNames };

src/renderer/wg_engine/tvgWgCompositor.cpp

@@ -87,6 +87,17 @@ WgPipelineBlendType WgCompositor::blendMethodToBlendType(BlendMethod blendMethod
 }
 
 
+RenderRegion WgCompositor::shrinkRenderRegion(RenderRegion& rect)
+{
+    // cut viewport to screen dimensions
+    int32_t xmin = std::max(0, std::min((int32_t)width, rect.x));
+    int32_t ymin = std::max(0, std::min((int32_t)height, rect.y));
+    int32_t xmax = std::max(0, std::min((int32_t)width, rect.x + rect.w));
+    int32_t ymax = std::max(0, std::min((int32_t)height, rect.y + rect.h));
+    return { xmin, ymin, xmax - xmin, ymax - ymin };
+}
+
+
 void WgCompositor::beginRenderPass(WGPUCommandEncoder commandEncoder, WgRenderStorage* target, bool clear)
 {
     assert(commandEncoder);
@@ -258,7 +269,19 @@ void WgCompositor::blendImage(WgContext& context, WgRenderDataPicture* renderDat
 };
 
 
-// TODO: use direct mask applience
+void WgCompositor::blendScene(WgContext& context, WgRenderStorage* src, WgCompose* cmp)
+{
+    assert(currentTarget);
+    RenderRegion rect = shrinkRenderRegion(cmp->aabb);
+    wgpuRenderPassEncoderSetScissorRect(renderPassEncoder, rect.x, rect.y, rect.w, rect.h);
+    wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 0);
+    wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 0, src->bindGroupTexure, 0, nullptr);
+    wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 1, bindGroupOpacities[cmp->opacity], 0, nullptr);
+    wgpuRenderPassEncoderSetPipeline(renderPassEncoder, pipelines->sceneBlend);
+    meshData.drawImage(context, renderPassEncoder);
+}
+
+
 void WgCompositor::composeShape(WgContext& context, WgRenderDataShape* renderData, WgRenderStorage* mask, CompositeMethod composeMethod)
 {
     assert(mask);
@@ -282,7 +305,6 @@ void WgCompositor::composeShape(WgContext& context, WgRenderDataShape* renderDat
 }
 
 
-// TODO: use direct mask applience
 void WgCompositor::composeStrokes(WgContext& context, WgRenderDataShape* renderData, WgRenderStorage* mask, CompositeMethod composeMethod)
 {
     assert(mask);
@@ -306,7 +328,6 @@ void WgCompositor::composeStrokes(WgContext& context, WgRenderDataShape* renderD
 }
 
 
-// TODO: use direct mask applience
 void WgCompositor::composeImage(WgContext& context, WgRenderDataPicture* renderData, WgRenderStorage* mask, CompositeMethod composeMethod)
 {
     assert(mask);
@@ -333,13 +354,7 @@ void WgCompositor::composeScene(WgContext& context, WgRenderStorage* src, WgRend
     assert(src);
     assert(mask);
     assert(renderPassEncoder);
-    // cut viewport to screen dimensions
+    RenderRegion rect = shrinkRenderRegion(cmp->aabb);
-    int32_t xmin = std::max(0, std::min((int32_t)width, cmp->aabb.x));
-    int32_t ymin = std::max(0, std::min((int32_t)height, cmp->aabb.y));
-    int32_t xmax = std::max(0, std::min((int32_t)width, cmp->aabb.x + cmp->aabb.w));
-    int32_t ymax = std::max(0, std::min((int32_t)height, cmp->aabb.y + cmp->aabb.h));
-    if ((xmin >= xmax) || (ymin >= ymax)) return;
-    RenderRegion rect { xmin, ymin, xmax - xmin, ymax - ymin };
     composeRegion(context, src, mask, cmp->method, rect);
 }
 
@@ -520,3 +535,19 @@ void WgCompositor::blend(WGPUCommandEncoder encoder, WgRenderStorage* src, WgRen
     wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, (width + 7) / 8, (height + 7) / 8, 1);
     wgpuComputePassEncoderEnd(computePassEncoder);
 }
+
+
+void WgCompositor::blit(WgContext& context, WGPUCommandEncoder encoder, WgRenderStorage* src, WGPUTextureView dstView) {
+    WGPURenderPassDepthStencilAttachment depthStencilAttachment{ .view = texViewStencil, .stencilLoadOp = WGPULoadOp_Load, .stencilStoreOp = WGPUStoreOp_Discard };
+    WGPURenderPassColorAttachment colorAttachment { .view = dstView, .loadOp = WGPULoadOp_Load, .storeOp = WGPUStoreOp_Store };
+    #ifdef __EMSCRIPTEN__
+    colorAttachment.depthSlice = WGPU_DEPTH_SLICE_UNDEFINED;
+    #endif
+    WGPURenderPassDescriptor renderPassDesc{ .colorAttachmentCount = 1, .colorAttachments = &colorAttachment, .depthStencilAttachment = &depthStencilAttachment };
+    WGPURenderPassEncoder renderPass = wgpuCommandEncoderBeginRenderPass(encoder, &renderPassDesc);
+    wgpuRenderPassEncoderSetBindGroup(renderPass, 0, src->bindGroupTexure, 0, nullptr);
+    wgpuRenderPassEncoderSetPipeline(renderPass, pipelines->blit);
+    meshData.drawImage(context, renderPass);
+    wgpuRenderPassEncoderEnd(renderPass);
+    wgpuRenderPassEncoderRelease(renderPass);
+}

src/renderer/wg_engine/tvgWgCompositor.h

@@ -58,6 +58,7 @@ private:
     static WgPipelineBlendType blendMethodToBlendType(BlendMethod blendMethod);
 
     void composeRegion(WgContext& context, WgRenderStorage* src, WgRenderStorage* mask, CompositeMethod composeMethod, RenderRegion& rect);
+    RenderRegion shrinkRenderRegion(RenderRegion& rect);
 public:
     // render target dimensions
     uint32_t width{};
@@ -77,6 +78,7 @@ public:
     void blendShape(WgContext& context, WgRenderDataShape* renderData, BlendMethod blendMethod);
     void blendStrokes(WgContext& context, WgRenderDataShape* renderData, BlendMethod blendMethod);
     void blendImage(WgContext& context, WgRenderDataPicture* renderData, BlendMethod blendMethod);
+    void blendScene(WgContext& context, WgRenderStorage* src, WgCompose* cmp);
 
     void composeShape(WgContext& context, WgRenderDataShape* renderData, WgRenderStorage* mask, CompositeMethod composeMethod);
     void composeStrokes(WgContext& context, WgRenderDataShape* renderData, WgRenderStorage* mask, CompositeMethod composeMethod);
@@ -90,6 +92,7 @@ public:
 
     void mergeMasks(WGPUCommandEncoder encoder, WgRenderStorage* mask0, WgRenderStorage* mask1);
     void blend(WGPUCommandEncoder encoder, WgRenderStorage* src, WgRenderStorage* dst, uint8_t opacity, BlendMethod blendMethod, WgRenderRasterType rasterType);
+    void blit(WgContext& context, WGPUCommandEncoder encoder, WgRenderStorage* src, WGPUTextureView dstView);
 };
 
 #endif // _TVG_WG_COMPOSITOR_H_

src/renderer/wg_engine/tvgWgPipelines.cpp

@@ -46,12 +46,12 @@ WGPURenderPipeline WgPipelines::createRenderPipeline(
     const WGPUShaderModule shaderModule, const char* vsEntryPoint, const char* fsEntryPoint,
     const WGPUPipelineLayout pipelineLayout,
     const WGPUVertexBufferLayout *vertexBufferLayouts, const uint32_t vertexBufferLayoutsCount,
-    const WGPUColorWriteMaskFlags writeMask,
+    const WGPUColorWriteMaskFlags writeMask, const WGPUTextureFormat colorTargetFormat,
     const WGPUCompareFunction stencilFunctionFrnt, const WGPUStencilOperation stencilOperationFrnt,
     const WGPUCompareFunction stencilFunctionBack, const WGPUStencilOperation stencilOperationBack,
     const WGPUPrimitiveState primitiveState, const WGPUMultisampleState multisampleState, const WGPUBlendState blendState)
 {
-    const WGPUColorTargetState colorTargetState { .format = WGPUTextureFormat_RGBA8Unorm, .blend = &blendState, .writeMask = writeMask };
+    const WGPUColorTargetState colorTargetState { .format = colorTargetFormat, .blend = &blendState, .writeMask = writeMask };
     const WGPUColorTargetState colorTargetStates[] { colorTargetState };
     const WGPUDepthStencilState depthStencilState {
         .format = WGPUTextureFormat_Stencil8, .depthCompare = WGPUCompareFunction_Always,
@@ -147,6 +147,7 @@ void WgPipelines::initialize(WgContext& context)
     const WGPUVertexBufferLayout vertexBufferLayoutsImage[] { vertexBufferLayoutPos, vertexBufferLayoutTex };
     const WGPUPrimitiveState primitiveState { .topology = WGPUPrimitiveTopology_TriangleList };
     const WGPUMultisampleState multisampleState { .count = 1, .mask = 0xFFFFFFFF, .alphaToCoverageEnabled = false };
+    const WGPUTextureFormat offscreenTargetFormat = WGPUTextureFormat_RGBA8Unorm;
 
     // blend states
     const WGPUBlendState blendStateSrc {
@@ -187,6 +188,13 @@ void WgPipelines::initialize(WgContext& context)
         layouts.layoutTexSampled,
         layouts.layoutTexSampled
     };
+    const WGPUBindGroupLayout bindGroupLayoutsSceneBlend[] = {
+        layouts.layoutTexSampled,
+        layouts.layoutBuffer1Un
+    };
+    const WGPUBindGroupLayout bindGroupLayoutsBlit[] = {
+        layouts.layoutTexSampled
+    };
     const WGPUBindGroupLayout bindGroupLayoutsMergeMasks[] = {
         layouts.layoutTexStrorage1RO,
         layouts.layoutTexStrorage1RO,
@@ -198,10 +206,6 @@ void WgPipelines::initialize(WgContext& context)
         layouts.layoutTexStrorage1WO,
         layouts.layoutBuffer1Un
     };
-    const WGPUBindGroupLayout bindGroupLayoutsCopy[] = {
-        layouts.layoutTexStrorage1RO,
-        layouts.layoutTexScreen1WO
-    };
 
     // pipeline layouts
     layoutStencil = createPipelineLayout(context.device, bindGroupLayoutsStencil, 2);
@@ -209,9 +213,10 @@ void WgPipelines::initialize(WgContext& context)
     layoutGradient = createPipelineLayout(context.device, bindGroupLayoutsGradient, 3);
     layoutImage = createPipelineLayout(context.device, bindGroupLayoutsImage, 3);
     layoutSceneComp = createPipelineLayout(context.device, bindGroupLayoutsSceneComp, 2);
+    layoutSceneBlend = createPipelineLayout(context.device, bindGroupLayoutsSceneBlend, 2);
+    layoutBlit = createPipelineLayout(context.device, bindGroupLayoutsBlit, 1);
     layoutBlend = createPipelineLayout(context.device, bindGroupLayoutsBlend, 4);
     layoutMergeMasks = createPipelineLayout(context.device, bindGroupLayoutsMergeMasks, 3);
-    layoutCopy = createPipelineLayout(context.device, bindGroupLayoutsCopy, 2);
 
     // graphics shader modules
     shaderStencil = createShaderModule(context.device, "The shader stencil", cShaderSrc_Stencil);
@@ -220,16 +225,17 @@ void WgPipelines::initialize(WgContext& context)
     shaderLinear = createShaderModule(context.device, "The shader linear", cShaderSrc_Linear);
     shaderImage = createShaderModule(context.device, "The shader image", cShaderSrc_Image);
     shaderSceneComp = createShaderModule(context.device, "The shader scene composition", cShaderSrc_Scene_Comp);
+    shaderSceneBlend = createShaderModule(context.device, "The shader scene blend", cShaderSrc_Scene_Blend);
+    shaderBlit = createShaderModule(context.device, "The shader blit", cShaderSrc_Blit);
     // computes shader modules
     shaderMergeMasks = createShaderModule(context.device, "The shader merge mask", cShaderSrc_MergeMasks);
-    shaderCopy = createShaderModule(context.device, "The shader copy", cShaderSrc_Copy);
 
     // render pipeline winding
     winding = createRenderPipeline(
         context.device, "The render pipeline winding",
         shaderStencil, "vs_main", "fs_main", layoutStencil,
         vertexBufferLayoutsShape, 1,
-        WGPUColorWriteMask_None,
+        WGPUColorWriteMask_None, offscreenTargetFormat,
         WGPUCompareFunction_Always, WGPUStencilOperation_IncrementWrap,
         WGPUCompareFunction_Always, WGPUStencilOperation_DecrementWrap,
         primitiveState, multisampleState, blendStateSrc);
@@ -238,7 +244,7 @@ void WgPipelines::initialize(WgContext& context)
         context.device, "The render pipeline even-odd",
         shaderStencil, "vs_main", "fs_main", layoutStencil,
         vertexBufferLayoutsShape, 1,
-        WGPUColorWriteMask_None,
+        WGPUColorWriteMask_None, offscreenTargetFormat,
         WGPUCompareFunction_Always, WGPUStencilOperation_Invert,
         WGPUCompareFunction_Always, WGPUStencilOperation_Invert,
         primitiveState, multisampleState, blendStateSrc);
@@ -247,7 +253,7 @@ void WgPipelines::initialize(WgContext& context)
         context.device, "The render pipeline direct",
         shaderStencil,"vs_main", "fs_main",  layoutStencil,
         vertexBufferLayoutsShape, 1,
-        WGPUColorWriteMask_None,
+        WGPUColorWriteMask_None, offscreenTargetFormat,
         WGPUCompareFunction_Always, WGPUStencilOperation_Replace,
         WGPUCompareFunction_Always, WGPUStencilOperation_Replace,
         primitiveState, multisampleState, blendStateSrc);
@@ -256,7 +262,7 @@ void WgPipelines::initialize(WgContext& context)
         context.device, "The render pipeline clip path",
         shaderStencil, "vs_main", "fs_main", layoutStencil,
         vertexBufferLayoutsShape, 1,
-        WGPUColorWriteMask_All,
+        WGPUColorWriteMask_All, offscreenTargetFormat,
         WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero,
         WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero,
         primitiveState, multisampleState, blendStateSrc);
@@ -267,7 +273,7 @@ void WgPipelines::initialize(WgContext& context)
             context.device, "The render pipeline solid",
             shaderSolid, "vs_main", "fs_main", layoutSolid,
             vertexBufferLayoutsShape, 1,
-            WGPUColorWriteMask_All,
+            WGPUColorWriteMask_All, offscreenTargetFormat,
             WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero,
             WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero,
             primitiveState, multisampleState, blendStates[i]);
@@ -279,7 +285,7 @@ void WgPipelines::initialize(WgContext& context)
             context.device, "The render pipeline radial",
             shaderRadial, "vs_main", "fs_main", layoutGradient,
             vertexBufferLayoutsShape, 1,
-            WGPUColorWriteMask_All,
+            WGPUColorWriteMask_All, offscreenTargetFormat,
             WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero,
             WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero,
             primitiveState, multisampleState, blendStates[i]);
@@ -291,7 +297,7 @@ void WgPipelines::initialize(WgContext& context)
             context.device, "The render pipeline linear",
             shaderLinear, "vs_main", "fs_main", layoutGradient,
             vertexBufferLayoutsShape, 1,
-            WGPUColorWriteMask_All,
+            WGPUColorWriteMask_All, offscreenTargetFormat,
             WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero,
             WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero,
             primitiveState, multisampleState, blendStates[i]);
@@ -303,13 +309,32 @@ void WgPipelines::initialize(WgContext& context)
             context.device, "The render pipeline image",
             shaderImage, "vs_main", "fs_main", layoutImage,
             vertexBufferLayoutsImage, 2,
-            WGPUColorWriteMask_All,
+            WGPUColorWriteMask_All, offscreenTargetFormat,
             WGPUCompareFunction_Always, WGPUStencilOperation_Zero,
             WGPUCompareFunction_Always, WGPUStencilOperation_Zero,
             primitiveState, multisampleState, blendStates[i]);
     }
 
-    // compute shader names
+    // render pipeline blit
+    blit = createRenderPipeline(context.device, "The render pipeline blit",
+            shaderBlit, "vs_main", "fs_main", layoutBlit,
+            vertexBufferLayoutsImage, 2,
+            // must be preferred screen pixel format
+            WGPUColorWriteMask_All, context.preferredFormat,
+            WGPUCompareFunction_Always, WGPUStencilOperation_Zero,
+            WGPUCompareFunction_Always, WGPUStencilOperation_Zero,
+            primitiveState, multisampleState, blendStateSrc);
+
+    // render pipeline blit
+    sceneBlend = createRenderPipeline(context.device, "The render pipeline scene blend",
+            shaderSceneBlend, "vs_main", "fs_main", layoutSceneBlend,
+            vertexBufferLayoutsImage, 2,
+            WGPUColorWriteMask_All, offscreenTargetFormat,
+            WGPUCompareFunction_Always, WGPUStencilOperation_Zero,
+            WGPUCompareFunction_Always, WGPUStencilOperation_Zero,
+            primitiveState, multisampleState, blendStateNrm);
+
+    // compose shader names
     const char* shaderComposeNames[] {
         "fs_main_None",
         "fs_main_ClipPath",
@@ -352,7 +377,7 @@ void WgPipelines::initialize(WgContext& context)
             context.device, "The render pipeline scene composition",
             shaderSceneComp, "vs_main", shaderComposeNames[i], layoutSceneComp,
             vertexBufferLayoutsImage, 2,
-            WGPUColorWriteMask_All,
+            WGPUColorWriteMask_All, offscreenTargetFormat,
             WGPUCompareFunction_Always, WGPUStencilOperation_Zero,
             WGPUCompareFunction_Always, WGPUStencilOperation_Zero,
             primitiveState, multisampleState, composeBlends[i]);
@@ -360,7 +385,6 @@ void WgPipelines::initialize(WgContext& context)
 
     // compute pipelines
     mergeMasks = createComputePipeline(context.device, "The pipeline merge masks", shaderMergeMasks, "cs_main", layoutMergeMasks);
-    copy = createComputePipeline(context.device, "The pipeline copy", shaderCopy, "cs_main", layoutCopy);
 
     // compute shader blend names
     const char* shaderBlendNames[] {
@@ -399,7 +423,9 @@ void WgPipelines::initialize(WgContext& context)
 
 void WgPipelines::releaseGraphicHandles(WgContext& context)
 {
+    releaseRenderPipeline(blit);
     releaseRenderPipeline(clipPath);
+    releaseRenderPipeline(sceneBlend);
     size_t pipesSceneCompCnt = sizeof(sceneComp) / sizeof(sceneComp[0]);
     for (uint32_t i = 0; i < pipesSceneCompCnt; i++)
         releaseRenderPipeline(sceneComp[i]);
@@ -412,11 +438,15 @@ void WgPipelines::releaseGraphicHandles(WgContext& context)
     releaseRenderPipeline(direct);
     releaseRenderPipeline(evenodd);
     releaseRenderPipeline(winding);
+    releasePipelineLayout(layoutBlit);
+    releasePipelineLayout(layoutSceneBlend);
     releasePipelineLayout(layoutSceneComp);
     releasePipelineLayout(layoutImage);
     releasePipelineLayout(layoutGradient);
     releasePipelineLayout(layoutSolid);
     releasePipelineLayout(layoutStencil);
+    releaseShaderModule(shaderBlit);
+    releaseShaderModule(shaderSceneBlend);
     releaseShaderModule(shaderSceneComp);
     releaseShaderModule(shaderImage);
     releaseShaderModule(shaderLinear);
@@ -434,16 +464,13 @@ void WgPipelines::releaseComputeHandles(WgContext& context)
         releaseComputePipeline(blendSolid[i]);
         releaseComputePipeline(blendGradient[i]);
     }
-    releaseComputePipeline(copy);
     releaseComputePipeline(mergeMasks);
     releasePipelineLayout(layoutBlend);
     releasePipelineLayout(layoutMergeMasks);
-    releasePipelineLayout(layoutCopy);
     releaseShaderModule(shaderBlendImage);
     releaseShaderModule(shaderBlendGradient);
     releaseShaderModule(shaderBlendSolid);
     releaseShaderModule(shaderMergeMasks);
-    releaseShaderModule(shaderCopy);
 }
 
 void WgPipelines::release(WgContext& context)

src/renderer/wg_engine/tvgWgPipelines.h

@@ -36,12 +36,13 @@ private:
     WGPUShaderModule shaderLinear{};
     WGPUShaderModule shaderImage{};
     WGPUShaderModule shaderSceneComp{};
+    WGPUShaderModule shaderSceneBlend{};
+    WGPUShaderModule shaderBlit{};
     // compute pipeline shaders
     WGPUShaderModule shaderMergeMasks;
     WGPUShaderModule shaderBlendSolid;
     WGPUShaderModule shaderBlendGradient;
     WGPUShaderModule shaderBlendImage;
-    WGPUShaderModule shaderCopy;
 private:
     // graphics pipeline layouts
     WGPUPipelineLayout layoutStencil{};
@@ -49,10 +50,11 @@ private:
     WGPUPipelineLayout layoutGradient{};
     WGPUPipelineLayout layoutImage{};
     WGPUPipelineLayout layoutSceneComp{};
+    WGPUPipelineLayout layoutSceneBlend{};
+    WGPUPipelineLayout layoutBlit{};
     // compute pipeline layouts
     WGPUPipelineLayout layoutMergeMasks{};
     WGPUPipelineLayout layoutBlend{};
-    WGPUPipelineLayout layoutCopy{};
 public:
     // graphics pipeline
     WgBindGroupLayouts layouts;
@@ -64,13 +66,14 @@ public:
     WGPURenderPipeline linear[3]{};
     WGPURenderPipeline image[3]{};
     WGPURenderPipeline sceneComp[12];
+    WGPURenderPipeline sceneBlend;
+    WGPURenderPipeline blit{};
     WGPURenderPipeline clipPath{};
     // compute pipeline
     WGPUComputePipeline mergeMasks;
     WGPUComputePipeline blendSolid[14];
     WGPUComputePipeline blendGradient[14];
     WGPUComputePipeline blendImage[14];
-    WGPUComputePipeline copy;
 private:
     void releaseGraphicHandles(WgContext& context);
     void releaseComputeHandles(WgContext& context);
@@ -82,7 +85,7 @@ private:
         const WGPUShaderModule shaderModule, const char* vsEntryPoint, const char* fsEntryPoint,
         const WGPUPipelineLayout pipelineLayout,
         const WGPUVertexBufferLayout *vertexBufferLayouts, const uint32_t vertexBufferLayoutsCount,
-        const WGPUColorWriteMaskFlags writeMask,
+        const WGPUColorWriteMaskFlags writeMask, const WGPUTextureFormat colorTargetFormat,
         const WGPUCompareFunction stencilFunctionFrnt, const WGPUStencilOperation stencilOperationFrnt,
         const WGPUCompareFunction stencilFunctionBack, const WGPUStencilOperation stencilOperationBack,
         const WGPUPrimitiveState primitiveState, const WGPUMultisampleState multisampleState, const WGPUBlendState blendState);

src/renderer/wg_engine/tvgWgRenderer.cpp

@@ -48,9 +48,6 @@ void WgRenderer::initialize()
 void WgRenderer::release()
 {
     disposeObjects();
-    mContext.releaseTexture(mTexScreen);
-    mContext.releaseTextureView(mTexViewScreen);
-    mContext.pipelines->layouts.releaseBindGroup(mBindGroupScreen);
     mStorageRoot.release(mContext);
     mRenderStoragePool.release(mContext);
     mRenderDataShapePool.release(mContext);
@@ -252,25 +249,14 @@ bool WgRenderer::sync()
     // get current texture
     WGPUSurfaceTexture surfaceTexture{};
     wgpuSurfaceGetCurrentTexture(mContext.surface, &surfaceTexture);
+    WGPUTextureView dstView = mContext.createTextureView(surfaceTexture.texture);
 
     // create command encoder
     const WGPUCommandEncoderDescriptor commandEncoderDesc{};
     WGPUCommandEncoder commandEncoder = wgpuDeviceCreateCommandEncoder(mContext.device, &commandEncoderDesc);
 
-    // copy render to screen with conversion rgba to bgra (screen format)
+    // show root offscreen buffer
-    const WGPUComputePassDescriptor computePassDescriptor{};
+    mCompositor.blit(mContext, commandEncoder, &mStorageRoot, dstView);
-    WGPUComputePassEncoder computePassEncoder = wgpuCommandEncoderBeginComputePass(commandEncoder, &computePassDescriptor);
-    wgpuComputePassEncoderSetBindGroup(computePassEncoder, 0, mStorageRoot.bindGroupRead, 0, nullptr);
-    wgpuComputePassEncoderSetBindGroup(computePassEncoder, 1, mBindGroupScreen, 0, nullptr);
-    wgpuComputePassEncoderSetPipeline(computePassEncoder, mContext.pipelines->copy);
-    wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, (mStorageRoot.width + 7) / 8, (mStorageRoot.height + 7) / 8, 1);
-    wgpuComputePassEncoderEnd(computePassEncoder);
-
-    // copy dst storage to temporary read only storage
-    const WGPUImageCopyTexture texSrc { .texture = mTexScreen };
-    const WGPUImageCopyTexture texDst { .texture = surfaceTexture.texture };
-    const WGPUExtent3D copySize { .width = mTargetSurface.w, .height = mTargetSurface.h, .depthOrArrayLayers = 1 };
-    wgpuCommandEncoderCopyTextureToTexture(commandEncoder, &texSrc, &texDst, &copySize);
 
     // release command encoder
     const WGPUCommandBufferDescriptor commandBufferDesc{};
@@ -278,6 +264,7 @@ bool WgRenderer::sync()
     wgpuQueueSubmit(mContext.queue, 1, &commandsBuffer);
     wgpuCommandBufferRelease(commandsBuffer);
     wgpuCommandEncoderRelease(commandEncoder);
+    mContext.releaseTextureView(dstView);
 
     return true;
 }
@@ -295,8 +282,8 @@ bool WgRenderer::target(WGPUInstance instance, WGPUSurface surface, uint32_t w,
 
     WGPUSurfaceConfiguration surfaceConfiguration {
         .device = mContext.device,
-        .format = WGPUTextureFormat_BGRA8Unorm,
+        .format = mContext.preferredFormat,
-        .usage = WGPUTextureUsage_CopyDst,
+        .usage = WGPUTextureUsage_RenderAttachment,
         .width = w, .height = h,
         #ifdef __EMSCRIPTEN__
         .presentMode = WGPUPresentMode_Fifo,
@@ -309,10 +296,6 @@ bool WgRenderer::target(WGPUInstance instance, WGPUSurface surface, uint32_t w,
     mRenderStoragePool.initialize(mContext, w, h);
     mStorageRoot.initialize(mContext, w, h);
     mCompositor.initialize(mContext, w, h);
-    // screen buffer
-    mTexScreen = mContext.createTexStorage(w, h, WGPUTextureFormat_BGRA8Unorm);
-    mTexViewScreen = mContext.createTextureView(mTexScreen);
-    mBindGroupScreen = mContext.pipelines->layouts.createBindGroupScreen1WO(mTexViewScreen);
     return true;
 }
 
@@ -355,12 +338,21 @@ bool WgRenderer::endComposite(Compositor* cmp)
         WgRenderStorage* src = mRenderStorageStack.last();
         mRenderStorageStack.pop();
         WgRenderStorage* dst = mRenderStorageStack.last();
-        // apply blend
+        // apply normal blend
-        mCompositor.blend(mCommandEncoder, src, dst, comp->opacity, comp->blend, WgRenderRasterType::Image);
+        if (comp->blend == BlendMethod::Normal) {
+            // begin previous render pass
+            mCompositor.beginRenderPass(mCommandEncoder, dst, false);
+            // apply blend
+            mCompositor.blendScene(mContext, src, comp);
+        // apply custom blend
+        } else {
+            // apply custom blend
+            mCompositor.blend(mCommandEncoder, src, dst, comp->opacity, comp->blend, WgRenderRasterType::Image);
+            // begin previous render pass
+            mCompositor.beginRenderPass(mCommandEncoder, dst, false);
+        }
         // back render targets to the pool
         mRenderStoragePool.free(mContext, src);
-        // begin previous render pass
-        mCompositor.beginRenderPass(mCommandEncoder, mRenderStorageStack.last(), false);
     } else { // finish composition
         // get source, mask and destination render storages
         WgRenderStorage* src = mRenderStorageStack.last();

src/renderer/wg_engine/tvgWgRenderer.h

@@ -75,11 +75,6 @@ private:
     WgPipelines mPipelines;
     WgCompositor mCompositor;
     
-    // screen buffer
-    WGPUTexture mTexScreen{};
-    WGPUTextureView mTexViewScreen{};
-    WGPUBindGroup mBindGroupScreen{};
-
     Surface mTargetSurface;
     BlendMethod mBlendMethod{};
     RenderRegion mViewport{};

src/renderer/wg_engine/tvgWgShaderSrc.cpp

@@ -291,6 +291,58 @@ fn fs_main_DarkenMask(in: VertexOutput) -> @location(0) vec4f {
 };
 )";
 
+
+//************************************************************************
+// graphics shader source: scene blend
+//************************************************************************
+
+const char* cShaderSrc_Scene_Blend = R"(
+struct VertexInput { @location(0) position: vec2f, @location(1) texCoord: vec2f };
+struct VertexOutput { @builtin(position) position: vec4f, @location(0) texCoord: vec2f };
+
+@group(0) @binding(0) var uSamplerSrc : sampler;
+@group(0) @binding(1) var uTextureSrc : texture_2d<f32>;
+@group(1) @binding(0) var<uniform> So : f32;
+
+@vertex
+fn vs_main(in: VertexInput) -> VertexOutput {
+    var out: VertexOutput;
+    out.position = vec4f(in.position.xy, 0.0, 1.0);
+    out.texCoord = in.texCoord;
+    return out;
+}
+
+@fragment
+fn fs_main(in: VertexOutput) -> @location(0) vec4f {
+    return textureSample(uTextureSrc, uSamplerSrc, in.texCoord.xy) * So;
+};
+)";
+
+//************************************************************************
+// graphics shader source: texture blit
+//************************************************************************
+
+const char* cShaderSrc_Blit = R"(
+struct VertexInput { @location(0) position: vec2f, @location(1) texCoord: vec2f };
+struct VertexOutput { @builtin(position) position: vec4f, @location(0) texCoord: vec2f };
+
+@group(0) @binding(0) var uSamplerSrc : sampler;
+@group(0) @binding(1) var uTextureSrc : texture_2d<f32>;
+
+@vertex
+fn vs_main(in: VertexInput) -> VertexOutput {
+    var out: VertexOutput;
+    out.position = vec4f(in.position.xy, 0.0, 1.0);
+    out.texCoord = in.texCoord;
+    return out;
+}
+
+@fragment
+fn fs_main(in: VertexOutput) -> @location(0) vec4f {
+    return textureSample(uTextureSrc, uSamplerSrc, in.texCoord.xy);
+};
+)";
+
 //************************************************************************
 // compute shader source: merge clip path masks
 //************************************************************************
@@ -518,17 +570,3 @@ fn cs_main_SoftLight(@builtin(global_invocation_id) id: vec3u) {
     textureStore(imageTgt, id.xy, postProcess(d, vec4f(Rc, 1.0)));
 };
 )";
-
-//************************************************************************
-// compute shader source: copy
-//************************************************************************
-
-const char* cShaderSrc_Copy = R"(
-@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)
-fn cs_main(@builtin(global_invocation_id) id: vec3u) {
-    textureStore(imageDst, id.xy, textureLoad(imageSrc, id.xy));
-};
-)";

src/renderer/wg_engine/tvgWgShaderSrc.h

@@ -30,6 +30,8 @@ extern const char* cShaderSrc_Linear;
 extern const char* cShaderSrc_Radial;
 extern const char* cShaderSrc_Image;
 extern const char* cShaderSrc_Scene_Comp;
+extern const char* cShaderSrc_Scene_Blend;
+extern const char* cShaderSrc_Blit;
 
 // compute shader sources: blend, compose and merge path
 extern const char* cShaderSrc_MergeMasks;
@@ -37,6 +39,5 @@ extern const char* cShaderSrc_BlendHeader_Solid;
 extern const char* cShaderSrc_BlendHeader_Gradient;
 extern const char* cShaderSrc_BlendHeader_Image;
 extern const char* cShaderSrc_Blend_Funcs;
-extern const char* cShaderSrc_Copy;
 
 #endif // _TVG_WG_SHEDER_SRC_H_