wg_engine: Introduce tiny effect for webgpu renderer

Issue: https://github.com/thorvg/thorvg/issues/3054

src/renderer/wg_engine/tvgWgCompositor.cpp

@@ -880,5 +880,33 @@ bool WgCompositor::fillEffect(WgContext& context, WgRenderStorage* dst, const Re
     wgpuComputePassEncoderEnd(computePassEncoder);
     wgpuComputePassEncoderRelease(computePassEncoder);
 
+    return true;
+}
+
+
+bool WgCompositor::tintEffect(WgContext& context, WgRenderStorage* dst, const RenderEffectFill* params, const WgCompose* compose)
+{
+    assert(dst);
+    assert(params);
+    assert(params->rd);
+    assert(compose->rdViewport);
+    assert(!renderPassEncoder);
+
+    auto renderDataParams = (WgRenderDataEffectParams*)params->rd;
+    auto aabb = compose->aabb;
+    auto viewport = compose->rdViewport;
+
+    copyTexture(&storageTemp0, dst, aabb);
+    WGPUComputePassDescriptor computePassDesc{ .label = "Compute pass tint" };
+    WGPUComputePassEncoder computePassEncoder = wgpuCommandEncoderBeginComputePass(commandEncoder, &computePassDesc);
+    wgpuComputePassEncoderSetBindGroup(computePassEncoder, 0, bindGroupStorageTemp, 0, nullptr);
+    wgpuComputePassEncoderSetBindGroup(computePassEncoder, 1, dst->bindGroupWrite, 0, nullptr);
+    wgpuComputePassEncoderSetBindGroup(computePassEncoder, 2, renderDataParams->bindGroupParams, 0, nullptr);
+    wgpuComputePassEncoderSetBindGroup(computePassEncoder, 3, viewport->bindGroupViewport, 0, nullptr);
+    wgpuComputePassEncoderSetPipeline(computePassEncoder, pipelines.tint_effect);
+    wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, (aabb.w - 1) / 128 + 1, aabb.h, 1);
+    wgpuComputePassEncoderEnd(computePassEncoder);
+    wgpuComputePassEncoderRelease(computePassEncoder);
+
     return true;
 }

src/renderer/wg_engine/tvgWgCompositor.h

@@ -115,6 +115,7 @@ public:
     bool gaussianBlur(WgContext& context, WgRenderStorage* dst, const RenderEffectGaussianBlur* params, const WgCompose* compose);
     bool dropShadow(WgContext& context, WgRenderStorage* dst, const RenderEffectDropShadow* params, const WgCompose* compose);
     bool fillEffect(WgContext& context, WgRenderStorage* dst, const RenderEffectFill* params, const WgCompose* compose);
+    bool tintEffect(WgContext& context, WgRenderStorage* dst, const RenderEffectFill* params, const WgCompose* compose);
 };
 
 #endif // _TVG_WG_COMPOSITOR_H_

src/renderer/wg_engine/tvgWgPipelines.cpp

@@ -454,11 +454,13 @@ void WgPipelines::initialize(WgContext& context)
     gaussian_vert = createComputePipeline(context.device, "The compute pipeline gaussian blur vertical",   shader_gauss, "cs_main_vert", layout_gauss);
     dropshadow    = createComputePipeline(context.device, "The compute pipeline drop shadow blend", shader_effects, "cs_main_drop_shadow", layout_effects);
     fill_effect   = createComputePipeline(context.device, "The compute pipeline fill effect", shader_effects, "cs_main_fill", layout_effects);
+    tint_effect   = createComputePipeline(context.device, "The compute pipeline tint effect", shader_effects, "cs_main_tint", layout_effects);
 }
 
 void WgPipelines::releaseGraphicHandles(WgContext& context)
 {
     // pipeline effects
+    releaseComputePipeline(tint_effect);
     releaseComputePipeline(fill_effect);
     releaseComputePipeline(dropshadow);
     releaseComputePipeline(gaussian_vert);

src/renderer/wg_engine/tvgWgPipelines.h

@@ -102,6 +102,7 @@ public:
     WGPUComputePipeline gaussian_vert{};
     WGPUComputePipeline dropshadow{};
     WGPUComputePipeline fill_effect{}; 
+    WGPUComputePipeline tint_effect{}; 
 private:
     void releaseGraphicHandles(WgContext& context);
     WGPUShaderModule createShaderModule(WGPUDevice device, const char* label, const char* code);

src/renderer/wg_engine/tvgWgRenderData.cpp

@@ -624,6 +624,15 @@ void WgRenderDataEffectParams::update(WgContext& context, const RenderEffectFill
 }
 
 
+void WgRenderDataEffectParams::update(WgContext& context, const RenderEffectTint* tint)
+{
+    assert(tint);
+    WgShaderTypeEffectParams effectParams;
+    effectParams.update(tint);
+    update(context, effectParams);
+}
+
+
 void WgRenderDataEffectParams::release(WgContext& context)
 {
     context.releaseBuffer(bufferParams);

src/renderer/wg_engine/tvgWgRenderData.h

@@ -205,6 +205,7 @@ struct WgRenderDataEffectParams
     void update(WgContext& context, const RenderEffectGaussianBlur* gaussian, const Matrix& transform);
     void update(WgContext& context, const RenderEffectDropShadow* dropShadow, const Matrix& transform);
     void update(WgContext& context, const RenderEffectFill* fill);
+    void update(WgContext& context, const RenderEffectTint* tint);
     void release(WgContext& context);
 };
 

src/renderer/wg_engine/tvgWgRenderer.cpp

@@ -542,7 +542,7 @@ void WgRenderer::prepare(RenderEffect* effect, const Matrix& transform)
         }
         renderData->update(mContext, renderEffect, transform);
         effect->valid = true;
-    }
+    } else
     // prepare fill
     if (effect->type == SceneEffect::Fill) {
         auto renderEffect = (RenderEffectFill*)effect;
@@ -553,6 +553,17 @@ void WgRenderer::prepare(RenderEffect* effect, const Matrix& transform)
         }
         renderData->update(mContext, renderEffect);
         effect->valid = true;
+    } else
+    // prepare tint
+    if (effect->type == SceneEffect::Tint) {
+        auto renderEffect = (RenderEffectTint*)effect;
+        auto renderData = (WgRenderDataEffectParams*)renderEffect->rd;
+        if (!renderData) {
+            renderData = mRenderDataEffectParamsPool.allocate(mContext);
+            renderEffect->rd = renderData;
+        }
+        renderData->update(mContext, renderEffect);
+        effect->valid = true;
     }
 }
 
@@ -598,6 +609,7 @@ bool WgRenderer::render(RenderCompositor* cmp, const RenderEffect* effect, TVG_U
         case SceneEffect::GaussianBlur: return mCompositor.gaussianBlur(mContext, dst, (RenderEffectGaussianBlur*)effect, comp);
         case SceneEffect::DropShadow: return mCompositor.dropShadow(mContext, dst, (RenderEffectDropShadow*)effect, comp);
         case SceneEffect::Fill: return mCompositor.fillEffect(mContext, dst, (RenderEffectFill*)effect, comp);
+        case SceneEffect::Tint: return mCompositor.tintEffect(mContext, dst, (RenderEffectFill*)effect, comp);
         default: return false;
     }
     return false;

src/renderer/wg_engine/tvgWgShaderSrc.cpp

@@ -843,11 +843,29 @@ fn cs_main_fill(@builtin(global_invocation_id) gid: vec3u) {
     let vmax = vec2u(viewport.zw);
 
     // tex coord
-    let uid = gid.xy + vmin;
+    let uid = min(gid.xy + vmin, vmax);
 
     let orig = textureLoad(imageSrc, uid);
     let fill = settings[0];
     let color = fill * orig.a * fill.a;
     textureStore(imageTrg, uid.xy, color);
 }
+
+@compute @workgroup_size(128, 1)
+fn cs_main_tint(@builtin(global_invocation_id) gid: vec3u) {
+    // decode viewport and settings
+    let vmin = vec2u(viewport.xy);
+    let vmax = vec2u(viewport.zw);
+
+    // tex coord
+    let uid = min(gid.xy + vmin, vmax);
+
+    let orig = textureLoad(imageSrc, uid);
+    let luma: f32 = dot(orig.rgb, vec3f(0.2125, 0.7154, 0.0721));
+    let black = settings[0];
+    let white = settings[1];
+    let intens = settings[2].r;
+    let color = mix(orig, mix(white, black, luma), intens) * orig.a;
+    textureStore(imageTrg, uid.xy, color);
+}
 )";

src/renderer/wg_engine/tvgWgShaderTypes.cpp

@@ -247,3 +247,16 @@ void WgShaderTypeEffectParams::update(const RenderEffectFill* fill)
     params[2] = fill->color[2] / 255.0f;
     params[3] = fill->color[3] / 255.0f;
 }
+
+void WgShaderTypeEffectParams::update(const RenderEffectTint* tint) 
+{
+    params[0] = tint->black[0] / 255.0f;
+    params[1] = tint->black[1] / 255.0f;
+    params[2] = tint->black[2] / 255.0f;
+    params[3] = 0.0f;
+    params[4] = tint->white[0] / 255.0f;
+    params[5] = tint->white[1] / 255.0f;
+    params[6] = tint->white[2] / 255.0f;
+    params[7] = 0.0f;
+    params[8] = tint->intensity / 255.0f;
+}

src/renderer/wg_engine/tvgWgShaderTypes.h

@@ -75,6 +75,7 @@ struct WgShaderTypeEffectParams
     // gaussian blur: [0]: sigma, [1]: scale, [2]: kernel size
     // drop shadow:   [0]: sigma, [1]: scale, [2]: kernel size, [4..7]: color, [8, 9]: offset
     // fill:          [0..3]: color
+    // tint:          [0..2]: black, [4..6]: white, [8]: intensity
     float params[4+4+4]{}; // settings: array<vec4f, 3>;
     uint32_t extend{};     // gaussian blur extend
     Point offset{};        // drop shadow offset
@@ -82,6 +83,7 @@ struct WgShaderTypeEffectParams
     void update(const RenderEffectGaussianBlur* gaussian, const Matrix& transform);
     void update(const RenderEffectDropShadow* dropShadow, const Matrix& transform);
     void update(const RenderEffectFill* fill);
+    void update(const RenderEffectTint* tint);
 };
 
 #endif // _TVG_WG_SHADER_TYPES_H_