From f024c54b8758405976dd0f8547652c0cfe5d619e Mon Sep 17 00:00:00 2001 From: Sergii Liebodkin Date: Thu, 27 Feb 2025 09:26:09 +0000 Subject: [PATCH] wg_engine: Introduce tiny effect for webgpu renderer Issue: https://github.com/thorvg/thorvg/issues/3054 --- src/renderer/wg_engine/tvgWgCompositor.cpp | 28 +++++++++++++++++++++ src/renderer/wg_engine/tvgWgCompositor.h | 1 + src/renderer/wg_engine/tvgWgPipelines.cpp | 2 ++ src/renderer/wg_engine/tvgWgPipelines.h | 1 + src/renderer/wg_engine/tvgWgRenderData.cpp | 9 +++++++ src/renderer/wg_engine/tvgWgRenderData.h | 1 + src/renderer/wg_engine/tvgWgRenderer.cpp | 14 ++++++++++- src/renderer/wg_engine/tvgWgShaderSrc.cpp | 20 ++++++++++++++- src/renderer/wg_engine/tvgWgShaderTypes.cpp | 13 ++++++++++ src/renderer/wg_engine/tvgWgShaderTypes.h | 2 ++ 10 files changed, 89 insertions(+), 2 deletions(-) diff --git a/src/renderer/wg_engine/tvgWgCompositor.cpp b/src/renderer/wg_engine/tvgWgCompositor.cpp index a65b508f..ba712645 100755 --- a/src/renderer/wg_engine/tvgWgCompositor.cpp +++ b/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; } \ No newline at end of file diff --git a/src/renderer/wg_engine/tvgWgCompositor.h b/src/renderer/wg_engine/tvgWgCompositor.h index 210309d7..b2dcb776 100755 --- a/src/renderer/wg_engine/tvgWgCompositor.h +++ b/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_ diff --git a/src/renderer/wg_engine/tvgWgPipelines.cpp b/src/renderer/wg_engine/tvgWgPipelines.cpp index 659ffc15..ded394fc 100755 --- a/src/renderer/wg_engine/tvgWgPipelines.cpp +++ b/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); diff --git a/src/renderer/wg_engine/tvgWgPipelines.h b/src/renderer/wg_engine/tvgWgPipelines.h index 73139e41..a5f13d1c 100755 --- a/src/renderer/wg_engine/tvgWgPipelines.h +++ b/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); diff --git a/src/renderer/wg_engine/tvgWgRenderData.cpp b/src/renderer/wg_engine/tvgWgRenderData.cpp index 265a99f8..d325c0e7 100755 --- a/src/renderer/wg_engine/tvgWgRenderData.cpp +++ b/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); diff --git a/src/renderer/wg_engine/tvgWgRenderData.h b/src/renderer/wg_engine/tvgWgRenderData.h index 9a2ab81e..23fa0802 100755 --- a/src/renderer/wg_engine/tvgWgRenderData.h +++ b/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); }; diff --git a/src/renderer/wg_engine/tvgWgRenderer.cpp b/src/renderer/wg_engine/tvgWgRenderer.cpp index 72d78405..49b337c4 100755 --- a/src/renderer/wg_engine/tvgWgRenderer.cpp +++ b/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; diff --git a/src/renderer/wg_engine/tvgWgShaderSrc.cpp b/src/renderer/wg_engine/tvgWgShaderSrc.cpp index 6f8abdd3..077bb2be 100755 --- a/src/renderer/wg_engine/tvgWgShaderSrc.cpp +++ b/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); +} )"; \ No newline at end of file diff --git a/src/renderer/wg_engine/tvgWgShaderTypes.cpp b/src/renderer/wg_engine/tvgWgShaderTypes.cpp index 3e42c0bd..bfe00408 100755 --- a/src/renderer/wg_engine/tvgWgShaderTypes.cpp +++ b/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; +} \ No newline at end of file diff --git a/src/renderer/wg_engine/tvgWgShaderTypes.h b/src/renderer/wg_engine/tvgWgShaderTypes.h index e5aa969a..ddfed6b6 100755 --- a/src/renderer/wg_engine/tvgWgShaderTypes.h +++ b/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; 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_