From 64ed756c2c2217fe9d0fbd5ad69e3bac05a447df Mon Sep 17 00:00:00 2001 From: Sergii Liebodkin Date: Wed, 19 Feb 2025 12:26:32 +0000 Subject: [PATCH] wg_engine: Introduce drop shadow effect for webgpu renderer Issue: https://github.com/thorvg/thorvg/issues/3054 Supported color, angle, distance, sigma, quality params --- src/renderer/wg_engine/tvgWgCompositor.cpp | 123 +++++++++++++++----- src/renderer/wg_engine/tvgWgCompositor.h | 10 +- src/renderer/wg_engine/tvgWgPipelines.cpp | 16 ++- src/renderer/wg_engine/tvgWgPipelines.h | 3 + src/renderer/wg_engine/tvgWgRenderData.cpp | 75 ++++++++++++ src/renderer/wg_engine/tvgWgRenderData.h | 27 +++++ src/renderer/wg_engine/tvgWgRenderer.cpp | 46 ++++++-- src/renderer/wg_engine/tvgWgRenderer.h | 1 + src/renderer/wg_engine/tvgWgShaderSrc.cpp | 40 ++++--- src/renderer/wg_engine/tvgWgShaderSrc.h | 2 +- src/renderer/wg_engine/tvgWgShaderTypes.cpp | 43 ++++++- src/renderer/wg_engine/tvgWgShaderTypes.h | 11 ++ 12 files changed, 329 insertions(+), 68 deletions(-) diff --git a/src/renderer/wg_engine/tvgWgCompositor.cpp b/src/renderer/wg_engine/tvgWgCompositor.cpp index 71a5f751..5a52681b 100755 --- a/src/renderer/wg_engine/tvgWgCompositor.cpp +++ b/src/renderer/wg_engine/tvgWgCompositor.cpp @@ -80,8 +80,10 @@ void WgCompositor::releasePools(WgContext& context) void WgCompositor::resize(WgContext& context, uint32_t width, uint32_t height) { // release existig handles if ((this->width != width) || (this->height != height)) { + context.layouts.releaseBindGroup(bindGroupStorageTemp); // release intermediate render storages - storageDstCopy.release(context); + storageTemp1.release(context); + storageTemp0.release(context); // release global stencil buffer handles context.releaseTextureView(texViewDepthStencilMS); context.releaseTexture(texDepthStencilMS); @@ -106,7 +108,9 @@ void WgCompositor::resize(WgContext& context, uint32_t width, uint32_t height) { texDepthStencilMS = context.createTexAttachement(width, height, WGPUTextureFormat_Depth24PlusStencil8, 4); texViewDepthStencilMS = context.createTextureView(texDepthStencilMS); // initialize intermediate render storages - storageDstCopy.initialize(context, width, height); + storageTemp0.initialize(context, width, height); + storageTemp1.initialize(context, width, height); + bindGroupStorageTemp = context.layouts.createBindGroupStrorage2RO(storageTemp0.texView, storageTemp1.texView); } } @@ -122,6 +126,25 @@ RenderRegion WgCompositor::shrinkRenderRegion(RenderRegion& rect) } +void WgCompositor::copyTexture(const WgRenderStorage* dst, const WgRenderStorage* src) +{ + const RenderRegion region = { 0, 0, (int32_t)src->width, (int32_t)src->height }; + copyTexture(dst, src, region); +} + + +void WgCompositor::copyTexture(const WgRenderStorage* dst, const WgRenderStorage* src, const RenderRegion& region) +{ + assert(dst); + assert(src); + assert(commandEncoder); + const WGPUImageCopyTexture texSrc { .texture = src->texture, .origin = { .x = (uint32_t)region.x, .y = (uint32_t)region.y } }; + const WGPUImageCopyTexture texDst { .texture = dst->texture, .origin = { .x = (uint32_t)region.x, .y = (uint32_t)region.y } }; + const WGPUExtent3D copySize { .width = (uint32_t)region.w, .height = (uint32_t)region.h, .depthOrArrayLayers = 1 }; + wgpuCommandEncoderCopyTextureToTexture(commandEncoder, &texSrc, &texDst, ©Size); +} + + void WgCompositor::beginRenderPass(WGPUCommandEncoder commandEncoder, WgRenderStorage* target, bool clear, WGPUColor clearColor) { assert(commandEncoder); @@ -313,10 +336,7 @@ void WgCompositor::blendShape(WgContext& context, WgRenderDataShape* renderData, // copy current render target data to dst storage WgRenderStorage *target = currentTarget; endRenderPass(); - const WGPUImageCopyTexture texSrc { .texture = target->texture }; - const WGPUImageCopyTexture texDst { .texture = storageDstCopy.texture }; - const WGPUExtent3D copySize { .width = width, .height = height, .depthOrArrayLayers = 1 }; - wgpuCommandEncoderCopyTextureToTexture(commandEncoder, &texSrc, &texDst, ©Size); + copyTexture(&storageTemp0, target); beginRenderPass(commandEncoder, target, false); // render shape with blend settings wgpuRenderPassEncoderSetScissorRect(renderPassEncoder, renderData->viewport.x, renderData->viewport.y, renderData->viewport.w, renderData->viewport.h); @@ -333,7 +353,7 @@ void WgCompositor::blendShape(WgContext& context, WgRenderDataShape* renderData, wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 0); wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 0, bindGroupViewMat, 0, nullptr); wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 1, renderData->bindGroupPaint, 0, nullptr); - wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 3, storageDstCopy.bindGroupTexure, 0, nullptr); + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 3, storageTemp0.bindGroupTexure, 0, nullptr); uint32_t blendMethodInd = (uint32_t)blendMethod; WgRenderSettings& settings = renderData->renderSettingsShape; if (settings.fillType == WgRenderSettingsType::Solid) { @@ -444,10 +464,7 @@ void WgCompositor::blendStrokes(WgContext& context, WgRenderDataShape* renderDat // copy current render target data to dst storage WgRenderStorage *target = currentTarget; endRenderPass(); - const WGPUImageCopyTexture texSrc { .texture = target->texture }; - const WGPUImageCopyTexture texDst { .texture = storageDstCopy.texture }; - const WGPUExtent3D copySize { .width = width, .height = height, .depthOrArrayLayers = 1 }; - wgpuCommandEncoderCopyTextureToTexture(commandEncoder, &texSrc, &texDst, ©Size); + copyTexture(&storageTemp0, target); beginRenderPass(commandEncoder, target, false); wgpuRenderPassEncoderSetScissorRect(renderPassEncoder, renderData->viewport.x, renderData->viewport.y, renderData->viewport.w, renderData->viewport.h); // draw strokes to stencil (first pass) @@ -463,7 +480,7 @@ void WgCompositor::blendStrokes(WgContext& context, WgRenderDataShape* renderDat wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 0); wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 0, bindGroupViewMat, 0, nullptr); wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 1, renderData->bindGroupPaint, 0, nullptr); - wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 3, storageDstCopy.bindGroupTexure, 0, nullptr); + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 3, storageTemp0.bindGroupTexure, 0, nullptr); uint32_t blendMethodInd = (uint32_t)blendMethod; WgRenderSettings& settings = renderData->renderSettingsStroke; if (settings.fillType == WgRenderSettingsType::Solid) { @@ -557,10 +574,7 @@ void WgCompositor::blendImage(WgContext& context, WgRenderDataPicture* renderDat // copy current render target data to dst storage WgRenderStorage *target = currentTarget; endRenderPass(); - const WGPUImageCopyTexture texSrc { .texture = target->texture }; - const WGPUImageCopyTexture texDst { .texture = storageDstCopy.texture }; - const WGPUExtent3D copySize { .width = width, .height = height, .depthOrArrayLayers = 1 }; - wgpuCommandEncoderCopyTextureToTexture(commandEncoder, &texSrc, &texDst, ©Size); + copyTexture(&storageTemp0, target); beginRenderPass(commandEncoder, target, false); // setup stencil rules wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 255); @@ -574,7 +588,7 @@ void WgCompositor::blendImage(WgContext& context, WgRenderDataPicture* renderDat wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 0, bindGroupViewMat, 0, nullptr); wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 1, renderData->bindGroupPaint, 0, nullptr); wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 2, renderData->bindGroupPicture, 0, nullptr); - wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 3, storageDstCopy.bindGroupTexure, 0, nullptr); + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 3, storageTemp0.bindGroupTexure, 0, nullptr); wgpuRenderPassEncoderSetPipeline(renderPassEncoder, pipelines.image_blend[blendMethodInd]); renderData->meshData.drawImage(context, renderPassEncoder); }; @@ -631,10 +645,7 @@ void WgCompositor::blendScene(WgContext& context, WgRenderStorage* scene, WgComp // copy current render target data to dst storage WgRenderStorage *target = currentTarget; endRenderPass(); - const WGPUImageCopyTexture texSrc { .texture = target->texture }; - const WGPUImageCopyTexture texDst { .texture = storageDstCopy.texture }; - const WGPUExtent3D copySize { .width = width, .height = height, .depthOrArrayLayers = 1 }; - wgpuCommandEncoderCopyTextureToTexture(commandEncoder, &texSrc, &texDst, ©Size); + copyTexture(&storageTemp0, target); beginRenderPass(commandEncoder, target, false); // blend scene uint32_t blendMethodInd = (uint32_t)compose->blend; @@ -642,7 +653,7 @@ void WgCompositor::blendScene(WgContext& context, WgRenderStorage* scene, WgComp wgpuRenderPassEncoderSetScissorRect(renderPassEncoder, rect.x, rect.y, rect.w, rect.h); wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 0); wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 0, scene->bindGroupTexure, 0, nullptr); - wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 1, storageDstCopy.bindGroupTexure, 0, nullptr); + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 1, storageTemp0.bindGroupTexure, 0, nullptr); wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 2, bindGroupOpacities[compose->opacity], 0, nullptr); wgpuRenderPassEncoderSetPipeline(renderPassEncoder, pipelines.scene_blend[blendMethodInd]); meshData.drawImage(context, renderPassEncoder); @@ -739,7 +750,7 @@ void WgCompositor::clearClipPath(WgContext& context, WgRenderDataPaint* paint) } -void WgCompositor::gaussianBlur(WgContext& context, WgRenderStorage* dst, const RenderEffectGaussianBlur* params, const WgCompose* compose) +bool WgCompositor::gaussianBlur(WgContext& context, WgRenderStorage* dst, const RenderEffectGaussianBlur* params, const WgCompose* compose) { assert(dst); assert(params); @@ -750,7 +761,7 @@ void WgCompositor::gaussianBlur(WgContext& context, WgRenderStorage* dst, const auto aabb = compose->aabb; auto viewport = compose->rdViewport; WgRenderStorage* sbuff = dst; - WgRenderStorage* dbuff = &storageDstCopy; + WgRenderStorage* dbuff = &storageTemp0; // begin compute pass WGPUComputePassDescriptor computePassDesc{ .label = "Compute pass gaussian blur" }; @@ -782,10 +793,64 @@ void WgCompositor::gaussianBlur(WgContext& context, WgRenderStorage* dst, const wgpuComputePassEncoderRelease(computePassEncoder); // if final result stored in intermidiate buffer we must copy result to destination buffer - if (sbuff == &storageDstCopy) { - const WGPUImageCopyTexture texSrc { .texture = sbuff->texture, .origin = { .x = (uint32_t)aabb.x, .y = (uint32_t)aabb.y } }; - const WGPUImageCopyTexture texDst { .texture = dbuff->texture, .origin = { .x = (uint32_t)aabb.x, .y = (uint32_t)aabb.y } }; - const WGPUExtent3D copySize { .width = (uint32_t)aabb.w, .height = (uint32_t)aabb.h, .depthOrArrayLayers = 1 }; - wgpuCommandEncoderCopyTextureToTexture(commandEncoder, &texSrc, &texDst, ©Size); + if (sbuff == &storageTemp0) + copyTexture(sbuff, dbuff, aabb); + + return true; +} + + +bool WgCompositor::dropShadow(WgContext& context, WgRenderStorage* dst, const RenderEffectDropShadow* params, const WgCompose* compose) +{ + assert(dst); + assert(params); + assert(params->rd); + assert(compose->rdViewport); + assert(!renderPassEncoder); + + auto renderDataDropShadow = (WgRenderDataDropShadow*)params->rd; + auto aabb = compose->aabb; + auto viewport = compose->rdViewport; + + { // apply blur + copyTexture(&storageTemp1, dst, aabb); + WgRenderStorage* sbuff = &storageTemp1; + WgRenderStorage* dbuff = &storageTemp0; + WGPUComputePassDescriptor computePassDesc{ .label = "Compute pass drop shadow blur" }; + WGPUComputePassEncoder computePassEncoder = wgpuCommandEncoderBeginComputePass(commandEncoder, &computePassDesc); + // horizontal blur + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 0, sbuff->bindGroupRead, 0, nullptr); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 1, dbuff->bindGroupWrite, 0, nullptr); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 2, renderDataDropShadow->bindGroupGaussian, 0, nullptr); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 3, viewport->bindGroupViewport, 0, nullptr); + wgpuComputePassEncoderSetPipeline(computePassEncoder, pipelines.gaussian_horz); + wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, (aabb.w - 1) / 128 + 1, aabb.h, 1); + std::swap(sbuff, dbuff); + // vertical blur + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 0, sbuff->bindGroupRead, 0, nullptr); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 1, dbuff->bindGroupWrite, 0, nullptr); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 2, renderDataDropShadow->bindGroupGaussian, 0, nullptr); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 3, viewport->bindGroupViewport, 0, nullptr); + wgpuComputePassEncoderSetPipeline(computePassEncoder, pipelines.gaussian_vert); + wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, aabb.w, (aabb.h - 1) / 128 + 1, 1); + std::swap(sbuff, dbuff); + wgpuComputePassEncoderEnd(computePassEncoder); + wgpuComputePassEncoderRelease(computePassEncoder); } + + { // blend origin (temp0), shadow (temp1) to destination + copyTexture(&storageTemp0, dst, aabb); + WGPUComputePassDescriptor computePassDesc{ .label = "Compute pass drop shadow blend" }; + WGPUComputePassEncoder computePassEncoder = wgpuCommandEncoderBeginComputePass(commandEncoder, &computePassDesc); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 0, bindGroupStorageTemp, 0, nullptr); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 1, dst->bindGroupWrite, 0, nullptr); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 2, renderDataDropShadow->bindGroupDropShadow, 0, nullptr); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 3, viewport->bindGroupViewport, 0, nullptr); + wgpuComputePassEncoderSetPipeline(computePassEncoder, pipelines.dropshadow); + 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 01234e59..0aa3778e 100755 --- a/src/renderer/wg_engine/tvgWgCompositor.h +++ b/src/renderer/wg_engine/tvgWgCompositor.h @@ -54,7 +54,9 @@ private: WGPUCommandEncoder commandEncoder{}; WgRenderStorage* currentTarget{}; // intermediate render storages - WgRenderStorage storageDstCopy; + WgRenderStorage storageTemp0; + WgRenderStorage storageTemp1; + WGPUBindGroup bindGroupStorageTemp{}; // composition and blend geometries WgMeshData meshData; // render target dimensions @@ -63,6 +65,9 @@ private: // viewport utilities RenderRegion shrinkRenderRegion(RenderRegion& rect); + void copyTexture(const WgRenderStorage* dst, const WgRenderStorage* src); + void copyTexture(const WgRenderStorage* dst, const WgRenderStorage* src, const RenderRegion& region); + // shapes void drawShape(WgContext& context, WgRenderDataShape* renderData); @@ -107,7 +112,8 @@ public: void blit(WgContext& context, WGPUCommandEncoder encoder, WgRenderStorage* src, WGPUTextureView dstView); // effects - void gaussianBlur(WgContext& context, WgRenderStorage* dst, const RenderEffectGaussianBlur* params, const WgCompose* compose); + bool gaussianBlur(WgContext& context, WgRenderStorage* dst, const RenderEffectGaussianBlur* params, const WgCompose* compose); + bool dropShadow(WgContext& context, WgRenderStorage* dst, const RenderEffectDropShadow* 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 367a3734..3b867d7b 100755 --- a/src/renderer/wg_engine/tvgWgPipelines.cpp +++ b/src/renderer/wg_engine/tvgWgPipelines.cpp @@ -185,6 +185,7 @@ void WgPipelines::initialize(WgContext& context) const WGPUBindGroupLayout bindGroupLayoutsBlit[] { layouts.layoutTexSampled }; // bind group layouts effects const WGPUBindGroupLayout bindGroupLayoutsGaussian[] { layouts.layoutTexStrorage1RO, layouts.layoutTexStrorage1WO, layouts.layoutBuffer1Un, layouts.layoutBuffer1Un }; + const WGPUBindGroupLayout bindGroupLayoutsDropShadow[] { layouts.layoutTexStrorage2RO, 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); @@ -222,6 +223,7 @@ void WgPipelines::initialize(WgContext& context) shader_blit = createShaderModule(context.device, "The shader blit", cShaderSrc_Blit); // shader effects shader_gaussian = createShaderModule(context.device, "The shader gaussian", cShaderSrc_GaussianBlur); + shader_dropshadow = createShaderModule(context.device, "The shader drop shadow", cShaderSrc_DropShadow); // layouts layout_stencil = createPipelineLayout(context.device, bindGroupLayoutsStencil, 2); @@ -242,6 +244,7 @@ void WgPipelines::initialize(WgContext& context) layout_blit = createPipelineLayout(context.device, bindGroupLayoutsBlit, 1); // layout effects layout_gaussian = createPipelineLayout(context.device, bindGroupLayoutsGaussian, 4); + layout_dropshadow = createPipelineLayout(context.device, bindGroupLayoutsDropShadow, 4); // render pipeline nonzero nonzero = createRenderPipeline( @@ -447,17 +450,16 @@ void WgPipelines::initialize(WgContext& context) depthStencilStateScene, multisampleStateX1); // compute pipeline gaussian blur - gaussian_horz = createComputePipeline( - context.device, "The compute pipeline gaussian blur horizontal", - shader_gaussian, "cs_main_horz", layout_gaussian); - gaussian_vert = createComputePipeline( - context.device, "The compute pipeline gaussian blur vertical", - shader_gaussian, "cs_main_vert", layout_gaussian); + gaussian_horz = createComputePipeline(context.device, "The compute pipeline gaussian blur horizontal", shader_gaussian, "cs_main_horz", layout_gaussian); + gaussian_vert = createComputePipeline(context.device, "The compute pipeline gaussian blur vertical", shader_gaussian, "cs_main_vert", layout_gaussian); + // compute pipeline drop shadow + dropshadow = createComputePipeline(context.device, "The compute pipeline drop shadow blend", shader_dropshadow, "cs_main", layout_dropshadow); } void WgPipelines::releaseGraphicHandles(WgContext& context) { // pipeline effects + releaseComputePipeline(dropshadow); releaseComputePipeline(gaussian_vert); releaseComputePipeline(gaussian_horz); // pipeline blit @@ -490,6 +492,7 @@ void WgPipelines::releaseGraphicHandles(WgContext& context) releaseRenderPipeline(evenodd); releaseRenderPipeline(nonzero); // layouts + releasePipelineLayout(layout_dropshadow); releasePipelineLayout(layout_gaussian); releasePipelineLayout(layout_blit); releasePipelineLayout(layout_scene_compose); @@ -504,6 +507,7 @@ void WgPipelines::releaseGraphicHandles(WgContext& context) releasePipelineLayout(layout_depth); releasePipelineLayout(layout_stencil); // shaders + releaseShaderModule(shader_dropshadow); releaseShaderModule(shader_gaussian); releaseShaderModule(shader_blit); releaseShaderModule(shader_scene_compose); diff --git a/src/renderer/wg_engine/tvgWgPipelines.h b/src/renderer/wg_engine/tvgWgPipelines.h index 86c7a9e4..215376b2 100755 --- a/src/renderer/wg_engine/tvgWgPipelines.h +++ b/src/renderer/wg_engine/tvgWgPipelines.h @@ -48,6 +48,7 @@ private: WGPUShaderModule shader_blit{}; // shader effects WGPUShaderModule shader_gaussian; + WGPUShaderModule shader_dropshadow; // layouts helpers WGPUPipelineLayout layout_stencil{}; @@ -68,6 +69,7 @@ private: WGPUPipelineLayout layout_blit{}; // layouts effects WGPUPipelineLayout layout_gaussian{}; + WGPUPipelineLayout layout_dropshadow{}; public: // pipelines stencil markup WGPURenderPipeline nonzero{}; @@ -98,6 +100,7 @@ public: // effects WGPUComputePipeline gaussian_horz{}; WGPUComputePipeline gaussian_vert{}; + WGPUComputePipeline dropshadow{}; 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 2fdda3d9..44f762d5 100755 --- a/src/renderer/wg_engine/tvgWgRenderData.cpp +++ b/src/renderer/wg_engine/tvgWgRenderData.cpp @@ -640,3 +640,78 @@ void WgRenderDataGaussianPool::release(WgContext& context) mPool.clear(); mList.clear(); } + +//*********************************************************************** +// WgRenderDataDropShadow +//*********************************************************************** + +void WgRenderDataDropShadow::update(WgContext& context, RenderEffectDropShadow* dropShadow, const Matrix& transform) +{ + assert(dropShadow); + // compute gaussian blur data + WgShaderTypeGaussianBlur gaussianSettings; + gaussianSettings.update(dropShadow, transform); + // update bind group and buffers + bool bufferGaussianChanged = context.allocateBufferUniform(bufferGaussian, &gaussianSettings.settings, sizeof(gaussianSettings.settings)); + if (bufferGaussianChanged) { + // update bind group + context.layouts.releaseBindGroup(bindGroupGaussian); + bindGroupGaussian = context.layouts.createBindGroupBuffer1Un(bufferGaussian); + } + // compute drop shadow data + WgShaderTypeDropShadow dropShadowSettings; + dropShadowSettings.update(dropShadow, transform); + // update bind group and buffers + bool bufferSettingsChanged = context.allocateBufferUniform(bufferSettings, &dropShadowSettings.settings, sizeof(dropShadowSettings.settings)); + if (bufferSettingsChanged) { + // update bind group + context.layouts.releaseBindGroup(bindGroupDropShadow); + bindGroupDropShadow = context.layouts.createBindGroupBuffer1Un(bufferSettings); + } + level = int(WG_GAUSSIAN_MAX_LEVEL * ((dropShadow->quality - 1) * 0.01f)) + 1; + extend = gaussianSettings.extend; + offset = dropShadowSettings.offset; +} + + +void WgRenderDataDropShadow::release(WgContext& context) +{ + context.releaseBuffer(bufferSettings); + context.releaseBuffer(bufferGaussian); + context.layouts.releaseBindGroup(bindGroupDropShadow); + context.layouts.releaseBindGroup(bindGroupGaussian); +} + +//*********************************************************************** +// WgRenderDataGaussianPool +//*********************************************************************** + +WgRenderDataDropShadow* WgRenderDataDropShadowPool::allocate(WgContext& context) +{ + WgRenderDataDropShadow* renderData{}; + if (mPool.count > 0) { + renderData = mPool.last(); + mPool.pop(); + } else { + renderData = new WgRenderDataDropShadow(); + mList.push(renderData); + } + return renderData; +} + + +void WgRenderDataDropShadowPool::free(WgContext& context, WgRenderDataDropShadow* renderData) +{ + if (renderData) mPool.push(renderData); +} + + +void WgRenderDataDropShadowPool::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 2d857b06..8f32aa93 100755 --- a/src/renderer/wg_engine/tvgWgRenderData.h +++ b/src/renderer/wg_engine/tvgWgRenderData.h @@ -215,4 +215,31 @@ public: void release(WgContext& context); }; +struct WgRenderDataDropShadow +{ + WGPUBindGroup bindGroupGaussian{}; + WGPUBindGroup bindGroupDropShadow{}; + WGPUBuffer bufferGaussian{}; + WGPUBuffer bufferSettings{}; + uint32_t extend{}; + uint32_t level{}; + Point offset{}; + + void update(WgContext& context, RenderEffectDropShadow* dropShadow, const Matrix& transform); + void release(WgContext& context); +}; + +class WgRenderDataDropShadowPool { +private: + // pool contains all created but unused render data for drop shadow + Array mPool; + // list contains all created render data for drop shadow + // to ensure that all created instances will be released + Array mList; +public: + WgRenderDataDropShadow* allocate(WgContext& context); + void free(WgContext& context, WgRenderDataDropShadow* 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 bb93e970..34e12580 100755 --- a/src/renderer/wg_engine/tvgWgRenderer.cpp +++ b/src/renderer/wg_engine/tvgWgRenderer.cpp @@ -53,6 +53,7 @@ void WgRenderer::release() mRenderDataShapePool.release(mContext); mRenderDataPicturePool.release(mContext); mRenderDataGaussianPool.release(mContext); + mRenderDataDropShadowPool.release(mContext); mRenderDataViewportPool.release(mContext); WgMeshDataPool::gMeshDataPool->release(mContext); @@ -531,12 +532,24 @@ void WgRenderer::prepare(RenderEffect* effect, const Matrix& transform) } renderDataGaussian->update(mContext, gaussianBlur, transform); effect->valid = true; + } else + // prepare drop shadow data + if (effect->type == SceneEffect::DropShadow) { + auto dropShadow = (RenderEffectDropShadow*)effect; + auto renderDataDropShadow = (WgRenderDataDropShadow*)dropShadow->rd; + if (!renderDataDropShadow) { + renderDataDropShadow = mRenderDataDropShadowPool.allocate(mContext); + dropShadow->rd = renderDataDropShadow; + } + renderDataDropShadow->update(mContext, dropShadow, transform); + effect->valid = true; } } bool WgRenderer::region(RenderEffect* effect) { + // update gaussian blur region if (effect->type == SceneEffect::GaussianBlur) { auto gaussian = (RenderEffectGaussianBlur*)effect; auto renderDataGaussian = (WgRenderDataGaussian*)gaussian->rd; @@ -549,6 +562,16 @@ bool WgRenderer::region(RenderEffect* effect) gaussian->extend.h = +renderDataGaussian->extend * 2; } return true; + } else + // update drop shadow region + if (effect->type == SceneEffect::DropShadow) { + auto dropShadow = (RenderEffectDropShadow*)effect; + auto renderDataDropShadow = (WgRenderDataDropShadow*)dropShadow->rd; + dropShadow->extend.x = -(renderDataDropShadow->extend + std::abs(renderDataDropShadow->offset.x)); + dropShadow->extend.w = +(renderDataDropShadow->extend + std::abs(renderDataDropShadow->offset.x)) * 2; + dropShadow->extend.y = -(renderDataDropShadow->extend + std::abs(renderDataDropShadow->offset.y)); + dropShadow->extend.h = +(renderDataDropShadow->extend + std::abs(renderDataDropShadow->offset.y)) * 2; + return true; } return false; } @@ -558,14 +581,13 @@ bool WgRenderer::render(RenderCompositor* cmp, const RenderEffect* effect, TVG_U { // we must to end current render pass to resolve ms texture before effect mCompositor.endRenderPass(); + WgCompose* comp = (WgCompose*)cmp; + WgRenderStorage* dst = mRenderStorageStack.last(); - // 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; + switch (effect->type) { + case SceneEffect::GaussianBlur: return mCompositor.gaussianBlur(mContext, dst, (RenderEffectGaussianBlur*)effect, comp); + case SceneEffect::DropShadow: return mCompositor.dropShadow(mContext, dst, (RenderEffectDropShadow*)effect, comp); + default: return false; } return false; } @@ -573,12 +595,12 @@ bool WgRenderer::render(RenderCompositor* cmp, const RenderEffect* effect, TVG_U void WgRenderer::dispose(RenderEffect* effect) { - // dispose gaussian blur data - if (effect->type == SceneEffect::GaussianBlur) { - auto gaussianBlur = (RenderEffectGaussianBlur*)effect; - mRenderDataGaussianPool.free(mContext, (WgRenderDataGaussian*)gaussianBlur->rd); - gaussianBlur->rd = nullptr; + switch (effect->type) { + case SceneEffect::GaussianBlur: mRenderDataGaussianPool.free(mContext, (WgRenderDataGaussian*)effect->rd); + case SceneEffect::DropShadow: mRenderDataDropShadowPool.free(mContext, (WgRenderDataDropShadow*)effect->rd); + default: effect->rd = nullptr; } + effect->rd = nullptr; }; diff --git a/src/renderer/wg_engine/tvgWgRenderer.h b/src/renderer/wg_engine/tvgWgRenderer.h index 4f870f61..ef7c3d7d 100755 --- a/src/renderer/wg_engine/tvgWgRenderer.h +++ b/src/renderer/wg_engine/tvgWgRenderer.h @@ -86,6 +86,7 @@ private: WgRenderDataShapePool mRenderDataShapePool; WgRenderDataPicturePool mRenderDataPicturePool; WgRenderDataGaussianPool mRenderDataGaussianPool; + WgRenderDataDropShadowPool mRenderDataDropShadowPool; WgRenderDataViewportPool mRenderDataViewportPool; // rendering context diff --git a/src/renderer/wg_engine/tvgWgShaderSrc.cpp b/src/renderer/wg_engine/tvgWgShaderSrc.cpp index 73b5c0d7..d6f6246d 100755 --- a/src/renderer/wg_engine/tvgWgShaderSrc.cpp +++ b/src/renderer/wg_engine/tvgWgShaderSrc.cpp @@ -708,19 +708,6 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f { // compute shader source: merge clip path masks //************************************************************************ -const char* cShaderSrc_MergeMasks = R"( -@group(0) @binding(0) var imageMsk0 : texture_storage_2d; -@group(1) @binding(0) var imageMsk1 : texture_storage_2d; -@group(2) @binding(0) var imageTrg : texture_storage_2d; - -@compute @workgroup_size(8, 8) -fn cs_main(@builtin(global_invocation_id) id: vec3u) { - let colorMsk0 = textureLoad(imageMsk0, id.xy); - let colorMsk1 = textureLoad(imageMsk1, id.xy); - textureStore(imageTrg, id.xy, colorMsk0 * colorMsk1); -} -)"; - const char* cShaderSrc_GaussianBlur = R"( @group(0) @binding(0) var imageSrc : texture_storage_2d; @group(1) @binding(0) var imageDst : texture_storage_2d; @@ -821,6 +808,31 @@ fn cs_main_vert(@builtin(global_invocation_id) gid: vec3u, // store result textureStore(imageDst, uid, color / sum); - //textureStore(imageDst, uid, vec4f(1.0, 0.0, 0.0, 1.0)); +} +)"; + +const char* cShaderSrc_DropShadow = R"( +@group(0) @binding(0) var imageSrc : texture_storage_2d; +@group(0) @binding(1) var imageSdw : texture_storage_2d; +@group(1) @binding(0) var imageTrg : texture_storage_2d; +@group(2) @binding(0) var settings: array; +@group(3) @binding(0) var viewport: vec4f; + +@compute @workgroup_size(128, 1) +fn cs_main(@builtin(global_invocation_id) gid: vec3u) { + // decode viewport and settings + let vmin = vec2u(viewport.xy); + let vmax = vec2u(viewport.zw); + let voff = vec2i(settings[1].xy); + + // tex coord + let uid = gid.xy + vmin; + let oid = clamp(vec2u(vec2i(uid) - voff), vmin, vmax); + + let orig = textureLoad(imageSrc, uid); + let blur = textureLoad(imageSdw, oid); + let shad = settings[0] * blur.a; + let color = orig + shad * (1.0 - orig.a); + textureStore(imageTrg, uid.xy, color); } )"; \ No newline at end of file diff --git a/src/renderer/wg_engine/tvgWgShaderSrc.h b/src/renderer/wg_engine/tvgWgShaderSrc.h index f445f5ed..126f5459 100755 --- a/src/renderer/wg_engine/tvgWgShaderSrc.h +++ b/src/renderer/wg_engine/tvgWgShaderSrc.h @@ -45,7 +45,7 @@ extern const char* cShaderSrc_Scene_Compose; extern const char* cShaderSrc_Blit; // compute shader sources: effects -extern const char* cShaderSrc_MergeMasks; extern const char* cShaderSrc_GaussianBlur; +extern const char* cShaderSrc_DropShadow; #endif // _TVG_WG_SHEDER_SRC_H_ diff --git a/src/renderer/wg_engine/tvgWgShaderTypes.cpp b/src/renderer/wg_engine/tvgWgShaderTypes.cpp index 9f214014..8f0bc2f5 100755 --- a/src/renderer/wg_engine/tvgWgShaderTypes.cpp +++ b/src/renderer/wg_engine/tvgWgShaderTypes.cpp @@ -200,14 +200,49 @@ void WgShaderTypeGradient::updateTexData(const Fill::ColorStop* stops, uint32_t // WgShaderTypeGaussianBlur //************************************************************************ -void WgShaderTypeGaussianBlur::update(const RenderEffectGaussianBlur* gaussian, const Matrix& transform) +void WgShaderTypeGaussianBlur::update(float sigma, const Matrix& transform) { - assert(gaussian); - const float sigma = gaussian->sigma; const float scale = std::sqrt(transform.e11 * transform.e11 + transform.e12 * transform.e12); const float kernel = std::min(WG_GAUSSIAN_KERNEL_SIZE_MAX, 2 * sigma * scale); // kernel size settings[0] = sigma; settings[1] = std::min(WG_GAUSSIAN_KERNEL_SIZE_MAX / kernel, scale); settings[2] = kernel; - extend = settings[2] * 2; + settings[3] = 0.0f; // unused + extend = settings[2] * 2; // kernel +} + + +void WgShaderTypeGaussianBlur::update(const RenderEffectGaussianBlur* gaussian, const Matrix& transform) +{ + assert(gaussian); + update(gaussian->sigma, transform); +} + + +void WgShaderTypeGaussianBlur::update(const RenderEffectDropShadow* dropShadow, const Matrix& transform) +{ + assert(dropShadow); + update(dropShadow->sigma, transform); +} + +//************************************************************************ +// WgShaderTypeDropShadow +//************************************************************************ + +void WgShaderTypeDropShadow::update(const RenderEffectDropShadow* dropShadow, const Matrix& transform) +{ + assert(dropShadow); + const float scale = std::sqrt(transform.e11 * transform.e11 + transform.e12 * transform.e12); + const float radian = tvg::deg2rad(90.0f - dropShadow->angle); + offset = {0, 0}; + if (dropShadow->distance > 0.0f) offset = { + +1.0f * dropShadow->distance * cosf(radian) * scale, + -1.0f * dropShadow->distance * sinf(radian) * scale + }; + settings[0] = dropShadow->color[0] / 255.0f; // red + settings[1] = dropShadow->color[1] / 255.0f; // green + settings[2] = dropShadow->color[2] / 255.0f; // blue + settings[3] = dropShadow->color[3] / 255.0f; // alpha + settings[4] = offset.x; + settings[5] = offset.y; } \ No newline at end of file diff --git a/src/renderer/wg_engine/tvgWgShaderTypes.h b/src/renderer/wg_engine/tvgWgShaderTypes.h index 85a64455..0bb3547c 100755 --- a/src/renderer/wg_engine/tvgWgShaderTypes.h +++ b/src/renderer/wg_engine/tvgWgShaderTypes.h @@ -74,7 +74,18 @@ struct WgShaderTypeGaussianBlur float settings[4]{}; // [0]: sigma, [1]: scale, [2]: kernel size, [3]: unused uint32_t extend{}; + void update(float sigma, const Matrix& transform); void update(const RenderEffectGaussianBlur* gaussian, const Matrix& transform); + void update(const RenderEffectDropShadow* dropShadow, const Matrix& transform); +}; + +// drop shadow settings: color, offset +struct WgShaderTypeDropShadow +{ + float settings[8]{}; // [0..3]: color, [4, 5]: offset + Point offset{}; + + void update(const RenderEffectDropShadow* dropShadow, const Matrix& transform); }; #endif // _TVG_WG_SHADER_TYPES_H_