From 857f1404e1a6430054072e5b4b6a0b623e8bb43e Mon Sep 17 00:00:00 2001 From: Sergii Liebodkin Date: Mon, 3 Feb 2025 12:57:57 +0000 Subject: [PATCH] wg_engine: gaussian blur basic implementation Introduce blur effect for webgpu renderer Issue: https://github.com/thorvg/thorvg/issues/3054 --- src/renderer/wg_engine/tvgWgCompositor.cpp | 65 ++++++++-- src/renderer/wg_engine/tvgWgCompositor.h | 4 + src/renderer/wg_engine/tvgWgPipelines.cpp | 21 ++++ src/renderer/wg_engine/tvgWgPipelines.h | 8 ++ src/renderer/wg_engine/tvgWgRenderData.cpp | 128 +++++++++++++++++++- src/renderer/wg_engine/tvgWgRenderData.h | 54 ++++++++- src/renderer/wg_engine/tvgWgRenderer.cpp | 90 +++++++++++--- src/renderer/wg_engine/tvgWgRenderer.h | 7 +- src/renderer/wg_engine/tvgWgShaderSrc.cpp | 94 +++++++++++++- src/renderer/wg_engine/tvgWgShaderSrc.h | 4 +- src/renderer/wg_engine/tvgWgShaderTypes.cpp | 23 ++++ src/renderer/wg_engine/tvgWgShaderTypes.h | 10 ++ 12 files changed, 474 insertions(+), 34 deletions(-) diff --git a/src/renderer/wg_engine/tvgWgCompositor.cpp b/src/renderer/wg_engine/tvgWgCompositor.cpp index b3f33646..20f2b88b 100755 --- a/src/renderer/wg_engine/tvgWgCompositor.cpp +++ b/src/renderer/wg_engine/tvgWgCompositor.cpp @@ -22,6 +22,7 @@ #include "tvgWgCompositor.h" #include "tvgWgShaderTypes.h" +#include void WgCompositor::initialize(WgContext& context, uint32_t width, uint32_t height) { @@ -115,8 +116,8 @@ 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)); + int32_t xmax = std::max(xmin, std::min((int32_t)width, rect.x + rect.w)); + int32_t ymax = std::max(ymin, std::min((int32_t)height, rect.y + rect.h)); return { xmin, ymin, xmax - xmin, ymax - ymin }; } @@ -150,11 +151,13 @@ void WgCompositor::beginRenderPass(WGPUCommandEncoder commandEncoder, WgRenderSt void WgCompositor::endRenderPass() { - assert(renderPassEncoder); - wgpuRenderPassEncoderEnd(renderPassEncoder); - wgpuRenderPassEncoderRelease(renderPassEncoder); - this->renderPassEncoder = nullptr; - this->currentTarget = nullptr; + if (currentTarget) { + assert(renderPassEncoder); + wgpuRenderPassEncoderEnd(renderPassEncoder); + wgpuRenderPassEncoderRelease(renderPassEncoder); + this->renderPassEncoder = nullptr; + this->currentTarget = nullptr; + } } @@ -733,4 +736,52 @@ void WgCompositor::clearClipPath(WgContext& context, WgRenderDataPaint* paint) wgpuRenderPassEncoderSetPipeline(renderPassEncoder, pipelines.clear_depth); renderData->meshDataBBox.drawFan(context, renderPassEncoder); } +} + + +void WgCompositor::gaussianBlur(WgContext& context, WgRenderStorage* dst, const RenderEffectGaussianBlur* params, const WgCompose* compose) +{ + assert(dst); + assert(params); + assert(params->rd); + assert(compose->rdViewport); + assert(!renderPassEncoder); + auto renderDataGaussian = (WgRenderDataGaussian*)params->rd; + auto viewport = compose->rdViewport; + for (uint32_t level = 0; level < renderDataGaussian->level; level++) { + // horizontal blur + if (params->direction != 2) { + const WGPUImageCopyTexture texSrc { .texture = dst->texture }; + const WGPUImageCopyTexture texDst { .texture = storageDstCopy.texture }; + const WGPUExtent3D copySize { .width = width, .height = height, .depthOrArrayLayers = 1 }; + wgpuCommandEncoderCopyTextureToTexture(commandEncoder, &texSrc, &texDst, ©Size); + WGPUComputePassDescriptor computePassDesc{ .label = "Compute pass gaussian blur horizontal" }; + WGPUComputePassEncoder computePassEncoder = wgpuCommandEncoderBeginComputePass(commandEncoder, &computePassDesc); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 0, storageDstCopy.bindGroupRead, 0, nullptr); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 1, dst->bindGroupWrite, 0, nullptr); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 2, renderDataGaussian->bindGroupGaussian, 0, nullptr); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 3, viewport->bindGroupViewport, 0, nullptr); + wgpuComputePassEncoderSetPipeline(computePassEncoder, pipelines.gaussian_horz); + wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, width / 16, height / 16, 1); + wgpuComputePassEncoderEnd(computePassEncoder); + wgpuComputePassEncoderRelease(computePassEncoder); + } + // vertical blur + if (params->direction != 1) { + const WGPUImageCopyTexture texSrc { .texture = dst->texture }; + const WGPUImageCopyTexture texDst { .texture = storageDstCopy.texture }; + const WGPUExtent3D copySize { .width = width, .height = height, .depthOrArrayLayers = 1 }; + wgpuCommandEncoderCopyTextureToTexture(commandEncoder, &texSrc, &texDst, ©Size); + WGPUComputePassDescriptor computePassDesc{ .label = "Compute pass gaussian blur vertical" }; + WGPUComputePassEncoder computePassEncoder = wgpuCommandEncoderBeginComputePass(commandEncoder, &computePassDesc); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 0, storageDstCopy.bindGroupRead, 0, nullptr); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 1, dst->bindGroupWrite, 0, nullptr); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 2, renderDataGaussian->bindGroupGaussian, 0, nullptr); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 3, viewport->bindGroupViewport, 0, nullptr); + wgpuComputePassEncoderSetPipeline(computePassEncoder, pipelines.gaussian_vert); + wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, width / 16, height / 16, 1); + wgpuComputePassEncoderEnd(computePassEncoder); + wgpuComputePassEncoderRelease(computePassEncoder); + } + } } \ No newline at end of file diff --git a/src/renderer/wg_engine/tvgWgCompositor.h b/src/renderer/wg_engine/tvgWgCompositor.h index 22505b85..01234e59 100755 --- a/src/renderer/wg_engine/tvgWgCompositor.h +++ b/src/renderer/wg_engine/tvgWgCompositor.h @@ -30,6 +30,7 @@ struct WgCompose: RenderCompositor { BlendMethod blend{}; RenderRegion aabb{}; + WgRenderDataViewport* rdViewport; }; class WgCompositor @@ -104,6 +105,9 @@ public: // blit render storage to texture view (f.e. screen buffer) void blit(WgContext& context, WGPUCommandEncoder encoder, WgRenderStorage* src, WGPUTextureView dstView); + + // effects + void gaussianBlur(WgContext& context, WgRenderStorage* dst, const RenderEffectGaussianBlur* 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 67c4a28a..9807b020 100755 --- a/src/renderer/wg_engine/tvgWgPipelines.cpp +++ b/src/renderer/wg_engine/tvgWgPipelines.cpp @@ -183,6 +183,8 @@ void WgPipelines::initialize(WgContext& context) const WGPUBindGroupLayout bindGroupLayoutsSceneCompose[] { layouts.layoutTexSampled, layouts.layoutTexSampled }; // bind group layouts blit const WGPUBindGroupLayout bindGroupLayoutsBlit[] { layouts.layoutTexSampled }; + // bind group layouts effects + const WGPUBindGroupLayout bindGroupLayoutsGaussian[] { layouts.layoutTexStrorage1RO, layouts.layoutTexStrorage1WO, layouts.layoutBuffer1Un, layouts.layoutBuffer1Un }; // depth stencil state markup const WGPUDepthStencilState depthStencilStateNonZero = makeDepthStencilState(WGPUCompareFunction_Always, false, WGPUCompareFunction_Always, WGPUStencilOperation_IncrementWrap, WGPUCompareFunction_Always, WGPUStencilOperation_DecrementWrap); @@ -218,6 +220,9 @@ void WgPipelines::initialize(WgContext& context) shader_scene_compose = createShaderModule(context.device, "The shader scene composition", cShaderSrc_Scene_Compose); // shader blit shader_blit = createShaderModule(context.device, "The shader blit", cShaderSrc_Blit); + // shader effects + shader_gaussian_horz = createShaderModule(context.device, "The shader gaussian horizontal", cShaderSrc_GaussianBlur_Horz); + shader_gaussian_vert = createShaderModule(context.device, "The shader gaussian vertical", cShaderSrc_GaussianBlur_Vert); // layouts layout_stencil = createPipelineLayout(context.device, bindGroupLayoutsStencil, 2); @@ -236,6 +241,8 @@ void WgPipelines::initialize(WgContext& context) layout_scene_compose = createPipelineLayout(context.device, bindGroupLayoutsSceneCompose, 2); // layout blit layout_blit = createPipelineLayout(context.device, bindGroupLayoutsBlit, 1); + // layout effects + layout_gaussian = createPipelineLayout(context.device, bindGroupLayoutsGaussian, 4); // render pipeline nonzero nonzero = createRenderPipeline( @@ -439,10 +446,21 @@ void WgPipelines::initialize(WgContext& context) layout_blit, vertexBufferLayoutsImage, 2, WGPUColorWriteMask_All, context.preferredFormat, blendStateSrc, // must be preferred screen pixel format depthStencilStateScene, multisampleStateX1); + + // compute pipeline gaussian blur + gaussian_horz = createComputePipeline( + context.device, "The compute pipeline gaussian blur horizontal", + shader_gaussian_horz, "cs_main", layout_gaussian); + gaussian_vert = createComputePipeline( + context.device, "The compute pipeline gaussian blur vertical", + shader_gaussian_vert, "cs_main", layout_gaussian); } void WgPipelines::releaseGraphicHandles(WgContext& context) { + // pipeline effects + releaseComputePipeline(gaussian_vert); + releaseComputePipeline(gaussian_horz); // pipeline blit releaseRenderPipeline(blit); // pipelines compose @@ -473,6 +491,7 @@ void WgPipelines::releaseGraphicHandles(WgContext& context) releaseRenderPipeline(evenodd); releaseRenderPipeline(nonzero); // layouts + releasePipelineLayout(layout_gaussian); releasePipelineLayout(layout_blit); releasePipelineLayout(layout_scene_compose); releasePipelineLayout(layout_scene_blend); @@ -486,6 +505,8 @@ void WgPipelines::releaseGraphicHandles(WgContext& context) releasePipelineLayout(layout_depth); releasePipelineLayout(layout_stencil); // shaders + releaseShaderModule(shader_gaussian_horz); + releaseShaderModule(shader_gaussian_vert); releaseShaderModule(shader_blit); releaseShaderModule(shader_scene_compose); releaseShaderModule(shader_scene_blend); diff --git a/src/renderer/wg_engine/tvgWgPipelines.h b/src/renderer/wg_engine/tvgWgPipelines.h index 98f39c9b..987aa787 100755 --- a/src/renderer/wg_engine/tvgWgPipelines.h +++ b/src/renderer/wg_engine/tvgWgPipelines.h @@ -46,6 +46,9 @@ private: WGPUShaderModule shader_scene_compose{}; // shader blit WGPUShaderModule shader_blit{}; + // shader effects + WGPUShaderModule shader_gaussian_horz{}; + WGPUShaderModule shader_gaussian_vert{}; // layouts helpers WGPUPipelineLayout layout_stencil{}; @@ -64,6 +67,8 @@ private: WGPUPipelineLayout layout_scene_compose{}; // layouts blit WGPUPipelineLayout layout_blit{}; + // layouts effects + WGPUPipelineLayout layout_gaussian{}; public: // pipelines stencil markup WGPURenderPipeline nonzero{}; @@ -91,6 +96,9 @@ public: WGPURenderPipeline scene_compose[11]{}; // pipeline blit WGPURenderPipeline blit{}; + // effects + WGPUComputePipeline gaussian_horz{}; + WGPUComputePipeline gaussian_vert{}; 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 eff3e9b3..5eb012e4 100755 --- a/src/renderer/wg_engine/tvgWgRenderData.cpp +++ b/src/renderer/wg_engine/tvgWgRenderData.cpp @@ -369,10 +369,10 @@ void WgRenderDataShape::updateAABB(const Matrix& tr) { Point p1 = Point{pMax.x, pMin.y} * tr; Point p2 = Point{pMin.x, pMax.y} * tr; Point p3 = Point{pMax.x, pMax.y} * tr; - aabb.x = std::min({p0.x, p1.x, p2.x, p3.x}); - aabb.y = std::min({p0.y, p1.y, p2.y, p3.y}); - aabb.w = std::max({p0.x, p1.x, p2.x, p3.x}) - aabb.x; - aabb.h = std::max({p0.y, p1.y, p2.y, p3.y}) - aabb.y; + aabb.pMin.x = std::min({p0.x, p1.x, p2.x, p3.x}); + aabb.pMin.y = std::min({p0.y, p1.y, p2.y, p3.y}); + aabb.pMax.x = std::max({p0.x, p1.x, p2.x, p3.x}); + aabb.pMax.y = std::max({p0.y, p1.y, p2.y, p3.y}); } @@ -448,7 +448,7 @@ void WgRenderDataShape::updateMeshes(WgContext& context, const RenderShape &rsha (this->meshGroupStrokesBBox.meshes.count > 0)) { updateAABB(tr); meshDataBBox.bbox(context, pMin, pMax); - } else aabb = {0, 0, 0, 0}; + } else aabb = {{0, 0}, {0, 0}}; } @@ -487,7 +487,7 @@ void WgRenderDataShape::releaseMeshes(WgContext& context) meshGroupShapes.release(context); pMin = {FLT_MAX, FLT_MAX}; pMax = {0.0f, 0.0f}; - aabb = {0, 0, 0, 0}; + aabb = {{0, 0}, {0, 0}}; clips.clear(); } @@ -600,3 +600,119 @@ void WgRenderDataPicturePool::release(WgContext& context) mPool.clear(); mList.clear(); } + +//*********************************************************************** +// WgRenderDataGaussian +//*********************************************************************** + +void WgRenderDataViewport::update(WgContext& context, const RenderRegion& region) { + WgShaderTypeVec4f viewport; + viewport.update(region); + bool bufferViewportChanged = context.allocateBufferUniform(bufferViewport, &viewport, sizeof(viewport)); + if (bufferViewportChanged) { + context.layouts.releaseBindGroup(bindGroupViewport); + bindGroupViewport = context.layouts.createBindGroupBuffer1Un(bufferViewport); + } +} + + +void WgRenderDataViewport::release(WgContext& context) { + context.releaseBuffer(bufferViewport); + context.layouts.releaseBindGroup(bindGroupViewport); +} + +//*********************************************************************** +// WgRenderDataViewportPool +//*********************************************************************** + +WgRenderDataViewport* WgRenderDataViewportPool::allocate(WgContext& context) +{ + WgRenderDataViewport* renderData{}; + if (mPool.count > 0) { + renderData = mPool.last(); + mPool.pop(); + } else { + renderData = new WgRenderDataViewport(); + mList.push(renderData); + } + return renderData; +} + + +void WgRenderDataViewportPool::free(WgContext& context, WgRenderDataViewport* renderData) +{ + if (renderData) mPool.push(renderData); +} + + +void WgRenderDataViewportPool::release(WgContext& context) +{ + ARRAY_FOREACH(p, mList) { + (*p)->release(context); + delete(*p); + } + mPool.clear(); + mList.clear(); +} + +//*********************************************************************** +// WgRenderDataGaussian +//*********************************************************************** + +void WgRenderDataGaussian::update(WgContext& context, RenderEffectGaussianBlur* gaussian, const Matrix& transform) +{ + assert(gaussian); + // compute gaussian blur data + WgShaderTypeGaussianBlur gaussianSettings; + gaussianSettings.update(gaussian, transform); + // update bind group and buffers + bool bufferSettingsChanged = context.allocateBufferUniform(bufferSettings, &gaussianSettings.settings, sizeof(gaussianSettings.settings)); + if (bufferSettingsChanged) { + // update bind group + context.layouts.releaseBindGroup(bindGroupGaussian); + bindGroupGaussian = context.layouts.createBindGroupBuffer1Un(bufferSettings); + } + level = int(WG_GAUSSIAN_MAX_LEVEL * ((gaussian->quality - 1) * 0.01f)) + 1; + extend = gaussianSettings.extend; +} + + +void WgRenderDataGaussian::release(WgContext& context) +{ + context.releaseBuffer(bufferSettings); + context.layouts.releaseBindGroup(bindGroupGaussian); +} + +//*********************************************************************** +// WgRenderDataGaussianPool +//*********************************************************************** + +WgRenderDataGaussian* WgRenderDataGaussianPool::allocate(WgContext& context) +{ + WgRenderDataGaussian* renderData{}; + if (mPool.count > 0) { + renderData = mPool.last(); + mPool.pop(); + } else { + renderData = new WgRenderDataGaussian(); + mList.push(renderData); + } + return renderData; +} + + +void WgRenderDataGaussianPool::free(WgContext& context, WgRenderDataGaussian* renderData) +{ + if (renderData) mPool.push(renderData); +} + + +void WgRenderDataGaussianPool::release(WgContext& context) +{ + ARRAY_FOREACH(p, mList) { + (*p)->release(context); + delete(*p); + } + mPool.clear(); + mList.clear(); +} diff --git a/src/renderer/wg_engine/tvgWgRenderData.h b/src/renderer/wg_engine/tvgWgRenderData.h index 528f2505..fdb1121f 100755 --- a/src/renderer/wg_engine/tvgWgRenderData.h +++ b/src/renderer/wg_engine/tvgWgRenderData.h @@ -26,6 +26,11 @@ #include "tvgWgPipelines.h" #include "tvgWgGeometry.h" +struct WgAabb { + Point pMin{}; + Point pMax{}; +}; + struct WgMeshData { WGPUBuffer bufferPosition{}; WGPUBuffer bufferTexCoord{}; @@ -99,7 +104,7 @@ struct WgRenderDataPaint WGPUBuffer bufferBlendSettings{}; WGPUBindGroup bindGroupPaint{}; RenderRegion viewport{}; - RenderRegion aabb{}; + WgAabb aabb{}; float opacity{}; Array clips; @@ -167,4 +172,51 @@ public: void release(WgContext& context); }; +struct WgRenderDataViewport +{ + WGPUBindGroup bindGroupViewport{}; + WGPUBuffer bufferViewport{}; + + void update(WgContext& context, const RenderRegion& region); + void release(WgContext& context); +}; + +class WgRenderDataViewportPool { +private: + // pool contains all created but unused render data for viewport + Array mPool; + // list contains all created render data for viewport + // to ensure that all created instances will be released + Array mList; +public: + WgRenderDataViewport* allocate(WgContext& context); + void free(WgContext& context, WgRenderDataViewport* renderData); + void release(WgContext& context); +}; + +#define WG_GAUSSIAN_MAX_LEVEL 3 +struct WgRenderDataGaussian +{ + WGPUBindGroup bindGroupGaussian{}; + WGPUBuffer bufferSettings{}; + uint32_t extend{}; + uint32_t level{}; + + void update(WgContext& context, RenderEffectGaussianBlur* gaussian, const Matrix& transform); + void release(WgContext& context); +}; + +class WgRenderDataGaussianPool { +private: + // pool contains all created but unused render data for gaussian filter + Array mPool; + // list contains all created render data for gaussian filter + // to ensure that all created instances will be released + Array mList; +public: + WgRenderDataGaussian* allocate(WgContext& context); + void free(WgContext& context, WgRenderDataGaussian* renderData); + void release(WgContext& context); +}; + #endif // _TVG_WG_RENDER_DATA_H_ diff --git a/src/renderer/wg_engine/tvgWgRenderer.cpp b/src/renderer/wg_engine/tvgWgRenderer.cpp index b2f75c30..678a01ef 100755 --- a/src/renderer/wg_engine/tvgWgRenderer.cpp +++ b/src/renderer/wg_engine/tvgWgRenderer.cpp @@ -49,6 +49,8 @@ void WgRenderer::release() // clear render data paint pools mRenderDataShapePool.release(mContext); mRenderDataPicturePool.release(mContext); + mRenderDataGaussianPool.release(mContext); + mRenderDataViewportPool.release(mContext); WgMeshDataPool::gMeshDataPool->release(mContext); // clear render storage pool @@ -182,6 +184,10 @@ bool WgRenderer::postRender() // pop root render storage to the render tree stack mRenderStorageStack.pop(); assert(mRenderStorageStack.count == 0); + // clear viewport list and store allocated handles to pool + ARRAY_FOREACH(p, mRenderDataViewportList) + mRenderDataViewportPool.free(mContext, *p); + mRenderDataViewportList.clear(); return true; } @@ -197,7 +203,14 @@ RenderRegion WgRenderer::region(RenderData data) { auto renderData = (WgRenderDataPaint*)data; if (renderData->type() == Type::Shape) { - return renderData->aabb; + Point v1 = renderData->aabb.pMin; + Point v2 = renderData->aabb.pMax; + RenderRegion renderRegion; + renderRegion.x = static_cast(nearbyint(v1.x)); + renderRegion.y = static_cast(nearbyint(v1.y)); + renderRegion.w = static_cast(nearbyint(v2.x)) - renderRegion.x; + renderRegion.h = static_cast(nearbyint(v2.y)) - renderRegion.y; + return renderRegion; } return { 0, 0, (int32_t)mTargetSurface.w, (int32_t)mTargetSurface.h }; } @@ -395,16 +408,24 @@ bool WgRenderer::surfaceConfigure(WGPUSurface surface, WgContext& context, uint3 RenderCompositor* WgRenderer::target(const RenderRegion& region, TVG_UNUSED ColorSpace cs, TVG_UNUSED CompositionFlag flags) { - mCompositorStack.push(new WgCompose); - mCompositorStack.last()->aabb = region; - return mCompositorStack.last(); + // create and setup compose data + WgCompose* compose = new WgCompose(); + compose->aabb = region; + if (flags & PostProcessing) { + compose->aabb = region; + compose->rdViewport = mRenderDataViewportPool.allocate(mContext); + compose->rdViewport->update(mContext, region); + mRenderDataViewportList.push(compose->rdViewport); + } + mCompositorStack.push(compose); + return compose; } bool WgRenderer::beginComposite(RenderCompositor* cmp, MaskMethod method, uint8_t opacity) { // save current composition settings - WgCompose* compose = (WgCompose *)cmp; + WgCompose* compose = (WgCompose*)cmp; compose->method = method; compose->opacity = opacity; compose->blend = mBlendMethod; @@ -424,8 +445,8 @@ bool WgRenderer::beginComposite(RenderCompositor* cmp, MaskMethod method, uint8_ bool WgRenderer::endComposite(RenderCompositor* cmp) { // get current composition settings - WgCompose* comp = (WgCompose *)cmp; - // end current render pass + WgCompose* comp = (WgCompose*)cmp; + // we must to end current render pass to run blend/composition mechanics mCompositor.endRenderPass(); // finish scene blending if (comp->method == MaskMethod::None) { @@ -463,30 +484,67 @@ bool WgRenderer::endComposite(RenderCompositor* cmp) } -void WgRenderer::prepare(TVG_UNUSED RenderEffect* effect, TVG_UNUSED const Matrix& transform) +void WgRenderer::prepare(RenderEffect* effect, const Matrix& transform) { - //TODO: prepare the effect + // prepare gaussian blur data + if (effect->type == SceneEffect::GaussianBlur) { + auto gaussianBlur = (RenderEffectGaussianBlur*)effect; + auto renderDataGaussian = (WgRenderDataGaussian*)gaussianBlur->rd; + if (!renderDataGaussian) { + renderDataGaussian = mRenderDataGaussianPool.allocate(mContext); + gaussianBlur->rd = renderDataGaussian; + } + renderDataGaussian->update(mContext, gaussianBlur, transform); + effect->valid = true; + } } -bool WgRenderer::region(TVG_UNUSED RenderEffect* effect) +bool WgRenderer::region(RenderEffect* effect) { - //TODO: Return if the current post effect requires the region expansion + if (effect->type == SceneEffect::GaussianBlur) { + auto gaussian = (RenderEffectGaussianBlur*)effect; + auto renderDataGaussian = (WgRenderDataGaussian*)gaussian->rd; + if (gaussian->direction != 2) { + gaussian->extend.x = -renderDataGaussian->extend; + gaussian->extend.w = +renderDataGaussian->extend * 2; + } + if (gaussian->direction != 1) { + gaussian->extend.y = -renderDataGaussian->extend; + gaussian->extend.h = +renderDataGaussian->extend * 2; + } + return true; + } return false; } -bool WgRenderer::render(TVG_UNUSED RenderCompositor* cmp, TVG_UNUSED const RenderEffect* effect, TVG_UNUSED bool direct) +bool WgRenderer::render(RenderCompositor* cmp, const RenderEffect* effect, TVG_UNUSED bool direct) { - TVGLOG("WG_ENGINE", "SceneEffect(%d) is not supported", (int)effect->type); + // we must to end current render pass to resolve ms texture before effect + mCompositor.endRenderPass(); + + // handle gaussian blur + if (effect->type == SceneEffect::GaussianBlur) { + WgCompose* comp = (WgCompose*)cmp; + WgRenderStorage* dst = mRenderStorageStack.last(); + RenderEffectGaussianBlur* gaussianBlur = (RenderEffectGaussianBlur*)effect; + mCompositor.gaussianBlur(mContext, dst, gaussianBlur, comp); + return true; + } return false; } -void WgRenderer::dispose(TVG_UNUSED RenderEffect* effect) +void WgRenderer::dispose(RenderEffect* effect) { - //TODO: dispose the effect -} + // dispose gaussian blur data + if (effect->type == SceneEffect::GaussianBlur) { + auto gaussianBlur = (RenderEffectGaussianBlur*)effect; + mRenderDataGaussianPool.free(mContext, (WgRenderDataGaussian*)gaussianBlur->rd); + gaussianBlur->rd = nullptr; + } +}; bool WgRenderer::preUpdate() diff --git a/src/renderer/wg_engine/tvgWgRenderer.h b/src/renderer/wg_engine/tvgWgRenderer.h index e40956d4..654e1956 100755 --- a/src/renderer/wg_engine/tvgWgRenderer.h +++ b/src/renderer/wg_engine/tvgWgRenderer.h @@ -53,10 +53,10 @@ public: bool beginComposite(RenderCompositor* cmp, MaskMethod method, uint8_t opacity) override; bool endComposite(RenderCompositor* cmp) override; - void prepare(TVG_UNUSED RenderEffect* effect, TVG_UNUSED const Matrix& transform) override; + void prepare(RenderEffect* effect, const Matrix& transform) override; bool region(RenderEffect* effect) override; bool render(RenderCompositor* cmp, const RenderEffect* effect, bool direct) override; - void dispose(TVG_UNUSED RenderEffect* effect) override; + void dispose(RenderEffect* effect) override; static WgRenderer* gen(); static bool init(uint32_t threads); @@ -77,6 +77,7 @@ private: WgRenderStorage mRenderStorageRoot; Array mCompositorStack; Array mRenderStorageStack; + Array mRenderDataViewportList; // render storage pool WgRenderStoragePool mRenderStoragePool; @@ -84,6 +85,8 @@ private: // render data paint pools WgRenderDataShapePool mRenderDataShapePool; WgRenderDataPicturePool mRenderDataPicturePool; + WgRenderDataGaussianPool mRenderDataGaussianPool; + WgRenderDataViewportPool mRenderDataViewportPool; // rendering context WgContext mContext; diff --git a/src/renderer/wg_engine/tvgWgShaderSrc.cpp b/src/renderer/wg_engine/tvgWgShaderSrc.cpp index b7f773af..501e7854 100755 --- a/src/renderer/wg_engine/tvgWgShaderSrc.cpp +++ b/src/renderer/wg_engine/tvgWgShaderSrc.cpp @@ -719,4 +719,96 @@ fn cs_main(@builtin(global_invocation_id) id: vec3u) { let colorMsk1 = textureLoad(imageMsk1, id.xy); textureStore(imageTrg, id.xy, colorMsk0 * colorMsk1); } -)"; \ No newline at end of file +)"; + +const char* cShaderSrc_GaussianBlur_Horz = R"( +@group(0) @binding(0) var imageSrc : texture_storage_2d; +@group(1) @binding(0) var imageDst : texture_storage_2d; +@group(2) @binding(0) var settings: vec4f; +@group(3) @binding(0) var viewport: vec4f; + +fn gaussian(x: f32, sigma: f32) -> f32 { + let a = 0.39894f / sigma; + let b = -(x * x) / (2.0 * sigma * sigma); + return a * exp(b); +} + +@compute @workgroup_size(16, 16) +fn cs_main(@builtin(global_invocation_id) id: vec3u) { + // id conversion + let iid = vec2i(id.xy); + // viewport decode + let xmin = i32(viewport.x); + let ymin = i32(viewport.y); + let xmax = i32(viewport.z); + let ymax = i32(viewport.w); + // settings decode + let sigma = settings.x; + let scale = settings.y; + let size = i32(settings.z); + + // draw borders points outside of viewport + if ((iid.x < xmin) || (iid.x > xmax) || (iid.y < ymin) || (iid.y > ymax)) { return; } + + // apply filter + var weight = gaussian(0.0, sigma); + var color = weight * textureLoad(imageSrc, id.xy); + var sum = weight; + for (var i: i32 = 1; i < size; i++) { + let ii = i32(f32(i) * scale); + let idneg = vec2i(clamp(iid.x - ii, xmin, xmax), iid.y); + let idpos = vec2i(clamp(iid.x + ii, xmin, xmax), iid.y); + weight = gaussian(f32(i) * scale, sigma); + color += (weight * textureLoad(imageSrc, vec2u(idneg))); + color += (weight * textureLoad(imageSrc, vec2u(idpos))); + sum += (2.0 * weight); + } + textureStore(imageDst, id.xy, color / sum); +} +)"; + +const char* cShaderSrc_GaussianBlur_Vert = R"( +@group(0) @binding(0) var imageSrc : texture_storage_2d; +@group(1) @binding(0) var imageDst : texture_storage_2d; +@group(2) @binding(0) var settings: vec4f; +@group(3) @binding(0) var viewport: vec4f; + +fn gaussian(x: f32, sigma: f32) -> f32 { + let a = 0.39894f / sigma; + let b = -(x * x) / (2.0 * sigma * sigma); + return a * exp(b); +} + +@compute @workgroup_size(16, 16) +fn cs_main(@builtin(global_invocation_id) id: vec3u) { + // id conversion + let iid = vec2i(id.xy); + // viewport decode + let xmin = i32(viewport.x); + let ymin = i32(viewport.y); + let xmax = i32(viewport.z); + let ymax = i32(viewport.w); + // settings decode + let sigma = settings.x; + let scale = settings.y; + let size = i32(settings.z); + + // draw borders points outside of viewport + if ((iid.x < xmin) || (iid.x > xmax) || (iid.y < ymin) || (iid.y > ymax)) { return; } + + // apply filter + var weight = gaussian(0.0, sigma); + var color = weight * textureLoad(imageSrc, id.xy); + var sum = weight; + for (var i: i32 = 1; i < size; i++) { + let ii = i32(f32(i) * scale); + let idneg = vec2i(iid.x, clamp(iid.y - ii, ymin, ymax)); + let idpos = vec2i(iid.x, clamp(iid.y + ii, ymin, ymax)); + weight = gaussian(f32(i) * scale, sigma); + color += (weight * textureLoad(imageSrc, vec2u(idneg))); + color += (weight * textureLoad(imageSrc, vec2u(idpos))); + sum += (2.0 * weight); + } + textureStore(imageDst, id.xy, color / sum); +} +)"; diff --git a/src/renderer/wg_engine/tvgWgShaderSrc.h b/src/renderer/wg_engine/tvgWgShaderSrc.h index 774ecd5a..68b7d6f8 100755 --- a/src/renderer/wg_engine/tvgWgShaderSrc.h +++ b/src/renderer/wg_engine/tvgWgShaderSrc.h @@ -44,7 +44,9 @@ extern const char* cShaderSrc_Scene_Compose; // shaders blit extern const char* cShaderSrc_Blit; -// compute shader sources: blend, compose and merge path +// compute shader sources: effects extern const char* cShaderSrc_MergeMasks; +extern const char* cShaderSrc_GaussianBlur_Vert; +extern const char* cShaderSrc_GaussianBlur_Horz; #endif // _TVG_WG_SHEDER_SRC_H_ diff --git a/src/renderer/wg_engine/tvgWgShaderTypes.cpp b/src/renderer/wg_engine/tvgWgShaderTypes.cpp index cae174e8..6c8b5604 100755 --- a/src/renderer/wg_engine/tvgWgShaderTypes.cpp +++ b/src/renderer/wg_engine/tvgWgShaderTypes.cpp @@ -116,6 +116,14 @@ void WgShaderTypeVec4f::update(const RenderColor& c) vec[3] = c.a / 255.0f; // alpha } +void WgShaderTypeVec4f::update(const RenderRegion& r) +{ + vec[0] = r.x; // left + vec[1] = r.y; // top + vec[2] = r.x + r.w - 1; // right + vec[3] = r.y + r.h - 1; // bottom +} + //************************************************************************ // WgShaderTypeGradient //************************************************************************ @@ -187,3 +195,18 @@ void WgShaderTypeGradient::updateTexData(const Fill::ColorStop* stops, uint32_t texData[ti * 4 + 3] = colorStopLast.a; } } + +//************************************************************************ +// WgShaderTypeGaussianBlur +//************************************************************************ + +void WgShaderTypeGaussianBlur::update(const RenderEffectGaussianBlur* gaussian, const Matrix& transform) +{ + assert(gaussian); + const float sigma = gaussian->sigma; + const float scale = std::sqrt(transform.e11 * transform.e11 + transform.e12 * transform.e12); + settings[0] = sigma; + settings[1] = scale; + settings[2] = 2 * sigma * scale; // kernel size + extend = settings[2] * 2; +} \ No newline at end of file diff --git a/src/renderer/wg_engine/tvgWgShaderTypes.h b/src/renderer/wg_engine/tvgWgShaderTypes.h index 7440a0af..8e27effb 100755 --- a/src/renderer/wg_engine/tvgWgShaderTypes.h +++ b/src/renderer/wg_engine/tvgWgShaderTypes.h @@ -52,6 +52,7 @@ struct WgShaderTypeVec4f WgShaderTypeVec4f(const RenderColor& c); void update(const ColorSpace colorSpace, uint8_t o); void update(const RenderColor& c); + void update(const RenderRegion& r); }; // sampler, texture, vec4f @@ -66,4 +67,13 @@ struct WgShaderTypeGradient void updateTexData(const Fill::ColorStop* stops, uint32_t stopCnt); }; +// gaussian settings: sigma, scale, extend +struct WgShaderTypeGaussianBlur +{ + float settings[4]{}; // [0]: sigma, [1]: scale, [2]: kernel size, [3]: unused + uint32_t extend{}; + + void update(const RenderEffectGaussianBlur* gaussian, const Matrix& transform); +}; + #endif // _TVG_WG_SHADER_TYPES_H_