wg_engine: blending and composition optimization

[issues 1479: Masking, InvMasking, LumaMasking, InvLumaMasking](#1479)

Computes composition and blending using simgle pass istead of two full screen passes

src/renderer/wg_engine/tvgWgBindGroups.cpp

@@ -35,6 +35,7 @@ WGPUBindGroupLayout WgBindGroupPicture::layout = nullptr;
 WGPUBindGroupLayout WgBindGroupTexture::layout = nullptr;
 WGPUBindGroupLayout WgBindGroupTextureStorage::layout = nullptr;
 WGPUBindGroupLayout WgBindGroupTextureSampled::layout = nullptr;
+WGPUBindGroupLayout WgBindGroupTexComposeBlend::layout = nullptr;
 WGPUBindGroupLayout WgBindGroupOpacity::layout = nullptr;
 WGPUBindGroupLayout WgBindGroupBlendMethod::layout = nullptr;
 WGPUBindGroupLayout WgBindGroupCompositeMethod::layout = nullptr;
@@ -386,6 +387,46 @@ void WgBindGroupTextureSampled::release()
 }
 
 
+WGPUBindGroupLayout WgBindGroupTexComposeBlend::getLayout(WGPUDevice device)
+{
+    if (layout) return layout;
+    const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] {
+        makeBindGroupLayoutEntryStorageTexture(0, WGPUStorageTextureAccess_ReadOnly),
+        makeBindGroupLayoutEntryStorageTexture(1, WGPUStorageTextureAccess_ReadOnly),
+        makeBindGroupLayoutEntryStorageTexture(2, WGPUStorageTextureAccess_ReadWrite)
+    };
+    layout = createBindGroupLayout(device, bindGroupLayoutEntries, 3);
+    assert(layout);
+    return layout;
+}
+
+
+void WgBindGroupTexComposeBlend::releaseLayout()
+{
+    releaseBindGroupLayout(layout);
+}
+
+
+void WgBindGroupTexComposeBlend::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexSrc, WGPUTextureView uTexMsk, WGPUTextureView uTexDst)
+{
+    release();
+    const WGPUBindGroupEntry bindGroupEntries[] {
+        makeBindGroupEntryTextureView(0, uTexSrc),
+        makeBindGroupEntryTextureView(1, uTexMsk),
+        makeBindGroupEntryTextureView(2, uTexDst)
+    };
+    mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 3);
+    assert(mBindGroup);
+}
+
+
+void WgBindGroupTexComposeBlend::release()
+{
+    releaseBindGroup(mBindGroup);
+}
+
+
+
 WGPUBindGroupLayout WgBindGroupOpacity::getLayout(WGPUDevice device)
 {
     if (layout) return layout;

src/renderer/wg_engine/tvgWgBindGroups.h

@@ -143,6 +143,21 @@ struct WgBindGroupTextureSampled : public WgBindGroup
     void release();
 };
 
+// @group(0)
+struct WgBindGroupTexComposeBlend : 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);
+    void release();
+};
+
+
 // @group(1 or 2)
 struct WgBindGroupOpacity : public WgBindGroup
 {

src/renderer/wg_engine/tvgWgPipelines.cpp

@@ -292,6 +292,28 @@ void WgPipelineCompose::initialize(WGPUDevice device)
 }
 
 
+void WgPipelineComposeBlend::initialize(WGPUDevice device)
+{
+    // bind groups and layouts
+    WGPUBindGroupLayout bindGroupLayouts[] = {
+        WgBindGroupTexComposeBlend::getLayout(device),
+        WgBindGroupCompositeMethod::getLayout(device),
+        WgBindGroupBlendMethod::getLayout(device),
+        WgBindGroupOpacity::getLayout(device)
+    };
+
+    // sheder source and labels
+    auto shaderSource = cShaderSource_PipelineComputeComposeBlend;
+    auto shaderLabel = "The compute shader compose blend";
+    auto pipelineLabel = "The compute pipeline compose blend";
+
+    // allocate all pipeline handles
+    allocate(device,
+             bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts),
+             shaderSource, shaderLabel, pipelineLabel);
+}
+
+
 void WgPipelineAntiAliasing::initialize(WGPUDevice device)
 {
     // bind groups and layouts
@@ -330,6 +352,7 @@ void WgPipelines::initialize(WgContext& context)
     computeClear.initialize(context.device);
     computeBlend.initialize(context.device);
     computeCompose.initialize(context.device);
+    computeComposeBlend.initialize(context.device);
     computeAntiAliasing.initialize(context.device);
     // store pipelines to context
     context.pipelines = this;
@@ -338,6 +361,7 @@ void WgPipelines::initialize(WgContext& context)
 
 void WgPipelines::release()
 {
+    WgBindGroupTexComposeBlend::layout = nullptr;
     WgBindGroupTextureSampled::releaseLayout();
     WgBindGroupTextureStorage::releaseLayout();
     WgBindGroupTexture::releaseLayout();
@@ -350,6 +374,7 @@ void WgPipelines::release()
     WgBindGroupCanvas::releaseLayout();
     // compute pipelines
     computeAntiAliasing.release();
+    computeComposeBlend.release();
     computeCompose.release();
     computeBlend.release();
     computeClear.release();

src/renderer/wg_engine/tvgWgPipelines.h

@@ -145,6 +145,20 @@ struct WgPipelineCompose: public WgComputePipeline
 };
 
 
+struct WgPipelineComposeBlend: public WgComputePipeline
+{
+    void initialize(WGPUDevice device) override;
+    void use(WGPUComputePassEncoder encoder, WgBindGroupTexComposeBlend& groupTexs, WgBindGroupCompositeMethod& groupComposeMethod, WgBindGroupBlendMethod& groupBlendMethod, WgBindGroupOpacity& groupOpacity)
+    {
+        set(encoder);
+        groupTexs.set(encoder, 0);
+        groupComposeMethod.set(encoder, 1);
+        groupBlendMethod.set(encoder, 2);
+        groupOpacity.set(encoder, 3);
+    }
+};
+
+
 struct WgPipelineAntiAliasing: public WgComputePipeline
 {
     void initialize(WGPUDevice device) override;
@@ -174,6 +188,7 @@ struct WgPipelines
     WgPipelineClear computeClear;
     WgPipelineBlend computeBlend;
     WgPipelineCompose computeCompose;
+    WgPipelineComposeBlend computeComposeBlend;
     WgPipelineAntiAliasing computeAntiAliasing;
 
     void initialize(WgContext& context);

src/renderer/wg_engine/tvgWgRenderTarget.cpp

@@ -185,6 +185,28 @@ void WgRenderStorage::compose(WGPUCommandEncoder commandEncoder, WgRenderStorage
 };
 
 
+void WgRenderStorage::composeBlend(
+    WgContext& context,
+    WGPUCommandEncoder commandEncoder,
+    WgRenderStorage* texSrc,
+    WgRenderStorage* texMsk,
+    WgBindGroupCompositeMethod* composeMethod,
+    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);
+    mPipelines->computeComposeBlend.use(computePassEncoder, composeBlend, *composeMethod, *blendMethod, *opacity);
+    dispatchWorkgroups(computePassEncoder);
+    endComputePass(computePassEncoder);
+    composeBlend.release();
+}
+
+
 void WgRenderStorage::antialias(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc)
 {
     assert(commandEncoder);

src/renderer/wg_engine/tvgWgRenderTarget.h

@@ -54,6 +54,14 @@ public:
     void clear(WGPUCommandEncoder commandEncoder);
     void blend(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc, WgBindGroupBlendMethod* blendMethod);
     void compose(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetMsk, WgBindGroupCompositeMethod* composeMethod, WgBindGroupOpacity* opacity);
+    void composeBlend(
+        WgContext& context,
+        WGPUCommandEncoder commandEncoder,
+        WgRenderStorage* texMsk,
+        WgRenderStorage* texSrc,
+        WgBindGroupCompositeMethod* composeMethod,
+        WgBindGroupBlendMethod* blendMethod,
+        WgBindGroupOpacity* opacity);
     void antialias(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc);
 private:
     void drawShape(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType);

src/renderer/wg_engine/tvgWgRenderer.cpp

@@ -392,20 +392,24 @@ bool WgRenderer::endComposite(TVG_UNUSED Compositor* cmp)
     } else {
         // end current render pass
         mRenderStorageStack.last()->endRenderPass();
-        // get two last render targets
+
+        // get source, mask and destination render storages
         WgRenderStorage* renderStorageSrc = mRenderStorageStack.last();
         mRenderStorageStack.pop();
         WgRenderStorage* renderStorageMsk = mRenderStorageStack.last();
         mRenderStorageStack.pop();
-    
-        // compose shape and mask
-        WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, cmp->opacity);
-        WgBindGroupCompositeMethod* composeMethod = mCompositeMethodPool.allocate(mContext, cmp->method);
-        renderStorageSrc->compose(mCommandEncoder, renderStorageMsk, composeMethod, opacity);
+        WgRenderStorage* renderStorageDst = mRenderStorageStack.last();
 
-        // blent scene to current render storage
+        // get compose, blend and opacity settings
+        WgBindGroupCompositeMethod* composeMethod = mCompositeMethodPool.allocate(mContext, cmp->method);
         WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod);
-        mRenderStorageStack.last()->blend(mCommandEncoder, renderStorageSrc, blendMethod);
+        WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, cmp->opacity);
+
+        // compose and blend
+        // dest = blend(dest, compose(src, msk, composeMethod), blendMethod, opacity)
+        renderStorageDst->composeBlend(mContext, mCommandEncoder,
+            renderStorageSrc, renderStorageMsk,
+            composeMethod, blendMethod, opacity);
 
         // back render targets to the pool
         mRenderStoragePool.free(mContext, renderStorageSrc);

src/renderer/wg_engine/tvgWgShaderSrc.cpp

@@ -453,6 +453,68 @@ fn cs_main( @builtin(global_invocation_id) id: vec3u) {
 }
 )";
 
+// pipeline shader modules compose blend
+const char* cShaderSource_PipelineComputeComposeBlend = R"(
+@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> composeMethod : u32;
+@group(2) @binding(0) var<uniform> blendMethod : u32;
+@group(3) @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 colorSrc = textureLoad(imageSrc, id.xy);
+    let colorMsk = textureLoad(imageMsk, id.xy);
+    let colorDst = textureLoad(imageDst, id.xy);
+
+    var color: vec3f = colorSrc.xyz;
+    var alpha: f32   = colorMsk.a;
+    let luma: vec3f  = vec3f(0.299, 0.587, 0.114);
+    switch composeMethod {
+        /* None           */ case 0u: { color = colorSrc.xyz; }
+        /* ClipPath       */ case 1u: { if (colorMsk.a == 0) { alpha = 0.0; }; }
+        /* AlphaMask      */ case 2u: { color = mix(colorMsk.xyz, colorSrc.xyz, colorSrc.a * colorMsk.b); }
+        /* InvAlphaMask   */ case 3u: { color = mix(colorSrc.xyz, colorMsk.xyz, colorSrc.a * colorMsk.b); alpha = 1.0 - colorMsk.b; }
+        /* LumaMask       */ case 4u: { color = colorSrc.xyz * dot(colorMsk.xyz, luma); }
+        /* InvLumaMask    */ case 5u: { color = colorSrc.xyz * (1.0 - dot(colorMsk.xyz, luma)); alpha = 1.0 - colorMsk.b; }
+        /* AddMask        */ case 6u: { color = colorSrc.xyz * colorSrc.a + colorMsk.xyz * (1.0 - colorSrc.a); }
+        /* SubtractMask   */ case 7u: { color = colorSrc.xyz * colorSrc.a - colorMsk.xyz * (1.0 - colorSrc.a); }
+        /* IntersectMask  */ case 8u: { color = colorSrc.xyz * min(colorSrc.a, colorMsk.a); }
+        /* DifferenceMask */ case 9u: { color = abs(colorMsk.xyz - colorSrc.xyz * (1.0 - colorMsk.a)); }
+        default: { color = colorSrc.xyz; }
+    }
+
+    let S: vec3f = color.xyz;
+    let D: vec3f = colorDst.xyz;
+    let Sa: f32  = alpha * opacity;
+    let Da: f32  = colorDst.a;
+    let One: vec3f = vec3(1.0);
+    switch blendMethod {
+        /* Normal     */ case 0u:  { color = (Sa * S) + (1.0 - Sa) * D; }
+        /* Add        */ case 1u:  { color = (S + D); }
+        /* Screen     */ case 2u:  { color = (S + D) - (S * D); }
+        /* Multiply   */ case 3u:  { color = (S * D); }
+        /* Overlay    */ case 4u:  { color = S * D; }
+        /* Difference */ case 5u:  { color = abs(S - D); }
+        /* Exclusion  */ case 6u:  { color = S + D - (2 * S * D); }
+        /* SrcOver    */ case 7u:  { color = S; }
+        /* Darken     */ case 8u:  { color = min(S, D); }
+        /* Lighten    */ case 9u:  { color = max(S, D); }
+        /* ColorDodge */ case 10u: { color = D / (One - S); }
+        /* ColorBurn  */ case 11u: { color = One - (One - D) / S; }
+        /* HardLight  */ case 12u: { color = (Sa * Da) - 2.0 * (Da - S) * (Sa - D); }
+        /* SoftLight  */ case 13u: { color = (One - 2 * S) * (D * D) + (2 * S * D); }
+        default:  { color = (Sa * S) + (1.0 - Sa) * D; }
+    }
+
+    textureStore(imageDst, id.xy, vec4f(color, Sa));
+}
+)";
+
 // pipeline shader modules anti-aliasing
 const char* cShaderSource_PipelineComputeAntiAlias = R"(
 @group(0) @binding(0) var imageSrc : texture_storage_2d<rgba8unorm, read_write>;

src/renderer/wg_engine/tvgWgShaderSrc.h

@@ -44,6 +44,7 @@ extern const char* cShaderSource_PipelineImage;
 extern const char* cShaderSource_PipelineComputeClear;
 extern const char* cShaderSource_PipelineComputeBlend;
 extern const char* cShaderSource_PipelineComputeCompose;
+extern const char* cShaderSource_PipelineComputeComposeBlend;
 extern const char* cShaderSource_PipelineComputeAntiAlias;
 
 #endif // _TVG_WG_SHADER_SRC_H_