diff --git a/src/renderer/wg_engine/tvgWgCompositor.cpp b/src/renderer/wg_engine/tvgWgCompositor.cpp index 5a52681b..a65b508f 100755 --- a/src/renderer/wg_engine/tvgWgCompositor.cpp +++ b/src/renderer/wg_engine/tvgWgCompositor.cpp @@ -757,7 +757,7 @@ bool WgCompositor::gaussianBlur(WgContext& context, WgRenderStorage* dst, const assert(params->rd); assert(compose->rdViewport); assert(!renderPassEncoder); - auto renderDataGaussian = (WgRenderDataGaussian*)params->rd; + auto renderData = (WgRenderDataEffectParams*)params->rd; auto aabb = compose->aabb; auto viewport = compose->rdViewport; WgRenderStorage* sbuff = dst; @@ -766,12 +766,12 @@ bool WgCompositor::gaussianBlur(WgContext& context, WgRenderStorage* dst, const // begin compute pass WGPUComputePassDescriptor computePassDesc{ .label = "Compute pass gaussian blur" }; WGPUComputePassEncoder computePassEncoder = wgpuCommandEncoderBeginComputePass(commandEncoder, &computePassDesc); - for (uint32_t level = 0; level < renderDataGaussian->level; level++) { + for (uint32_t level = 0; level < renderData->level; level++) { // horizontal blur if (params->direction != 2) { wgpuComputePassEncoderSetBindGroup(computePassEncoder, 0, sbuff->bindGroupRead, 0, nullptr); wgpuComputePassEncoderSetBindGroup(computePassEncoder, 1, dbuff->bindGroupWrite, 0, nullptr); - wgpuComputePassEncoderSetBindGroup(computePassEncoder, 2, renderDataGaussian->bindGroupGaussian, 0, nullptr); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 2, renderData->bindGroupParams, 0, nullptr); wgpuComputePassEncoderSetBindGroup(computePassEncoder, 3, viewport->bindGroupViewport, 0, nullptr); wgpuComputePassEncoderSetPipeline(computePassEncoder, pipelines.gaussian_horz); wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, (aabb.w - 1) / 128 + 1, aabb.h, 1); @@ -781,7 +781,7 @@ bool WgCompositor::gaussianBlur(WgContext& context, WgRenderStorage* dst, const if (params->direction != 1) { wgpuComputePassEncoderSetBindGroup(computePassEncoder, 0, sbuff->bindGroupRead, 0, nullptr); wgpuComputePassEncoderSetBindGroup(computePassEncoder, 1, dbuff->bindGroupWrite, 0, nullptr); - wgpuComputePassEncoderSetBindGroup(computePassEncoder, 2, renderDataGaussian->bindGroupGaussian, 0, nullptr); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 2, renderData->bindGroupParams, 0, nullptr); wgpuComputePassEncoderSetBindGroup(computePassEncoder, 3, viewport->bindGroupViewport, 0, nullptr); wgpuComputePassEncoderSetPipeline(computePassEncoder, pipelines.gaussian_vert); wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, aabb.w, (aabb.h - 1) / 128 + 1, 1); @@ -808,7 +808,7 @@ bool WgCompositor::dropShadow(WgContext& context, WgRenderStorage* dst, const Re assert(compose->rdViewport); assert(!renderPassEncoder); - auto renderDataDropShadow = (WgRenderDataDropShadow*)params->rd; + auto renderDataParams = (WgRenderDataEffectParams*)params->rd; auto aabb = compose->aabb; auto viewport = compose->rdViewport; @@ -821,7 +821,7 @@ bool WgCompositor::dropShadow(WgContext& context, WgRenderStorage* dst, const Re // 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, 2, renderDataParams->bindGroupParams, 0, nullptr); wgpuComputePassEncoderSetBindGroup(computePassEncoder, 3, viewport->bindGroupViewport, 0, nullptr); wgpuComputePassEncoderSetPipeline(computePassEncoder, pipelines.gaussian_horz); wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, (aabb.w - 1) / 128 + 1, aabb.h, 1); @@ -829,7 +829,7 @@ bool WgCompositor::dropShadow(WgContext& context, WgRenderStorage* dst, const Re // 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, 2, renderDataParams->bindGroupParams, 0, nullptr); wgpuComputePassEncoderSetBindGroup(computePassEncoder, 3, viewport->bindGroupViewport, 0, nullptr); wgpuComputePassEncoderSetPipeline(computePassEncoder, pipelines.gaussian_vert); wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, aabb.w, (aabb.h - 1) / 128 + 1, 1); @@ -844,7 +844,7 @@ bool WgCompositor::dropShadow(WgContext& context, WgRenderStorage* dst, const Re 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, 2, renderDataParams->bindGroupParams, 0, nullptr); wgpuComputePassEncoderSetBindGroup(computePassEncoder, 3, viewport->bindGroupViewport, 0, nullptr); wgpuComputePassEncoderSetPipeline(computePassEncoder, pipelines.dropshadow); wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, (aabb.w - 1) / 128 + 1, aabb.h, 1); @@ -852,5 +852,33 @@ bool WgCompositor::dropShadow(WgContext& context, WgRenderStorage* dst, const Re wgpuComputePassEncoderRelease(computePassEncoder); } + return true; +} + + +bool WgCompositor::fillEffect(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 fill" }; + 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.fill_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 0aa3778e..210309d7 100755 --- a/src/renderer/wg_engine/tvgWgCompositor.h +++ b/src/renderer/wg_engine/tvgWgCompositor.h @@ -114,6 +114,7 @@ public: // effects 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); }; #endif // _TVG_WG_COMPOSITOR_H_ diff --git a/src/renderer/wg_engine/tvgWgPipelines.cpp b/src/renderer/wg_engine/tvgWgPipelines.cpp index 3b867d7b..659ffc15 100755 --- a/src/renderer/wg_engine/tvgWgPipelines.cpp +++ b/src/renderer/wg_engine/tvgWgPipelines.cpp @@ -184,8 +184,8 @@ void WgPipelines::initialize(WgContext& context) // bind group layouts blit 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 }; + const WGPUBindGroupLayout bindGroupLayoutsGauss[] { layouts.layoutTexStrorage1RO, layouts.layoutTexStrorage1WO, layouts.layoutBuffer1Un, layouts.layoutBuffer1Un }; + const WGPUBindGroupLayout bindGroupLayoutsEffects[] { 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,8 +222,8 @@ void WgPipelines::initialize(WgContext& context) // shader blit 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); + shader_gauss = createShaderModule(context.device, "The shader effects", cShaderSrc_GaussianBlur); + shader_effects = createShaderModule(context.device, "The shader effects", cShaderSrc_Effects); // layouts layout_stencil = createPipelineLayout(context.device, bindGroupLayoutsStencil, 2); @@ -243,8 +243,8 @@ void WgPipelines::initialize(WgContext& context) // layout blit layout_blit = createPipelineLayout(context.device, bindGroupLayoutsBlit, 1); // layout effects - layout_gaussian = createPipelineLayout(context.device, bindGroupLayoutsGaussian, 4); - layout_dropshadow = createPipelineLayout(context.device, bindGroupLayoutsDropShadow, 4); + layout_gauss = createPipelineLayout(context.device, bindGroupLayoutsGauss, 4); + layout_effects = createPipelineLayout(context.device, bindGroupLayoutsEffects, 4); // render pipeline nonzero nonzero = createRenderPipeline( @@ -450,15 +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); - // compute pipeline drop shadow - dropshadow = createComputePipeline(context.device, "The compute pipeline drop shadow blend", shader_dropshadow, "cs_main", layout_dropshadow); + gaussian_horz = createComputePipeline(context.device, "The compute pipeline gaussian blur horizontal", shader_gauss, "cs_main_horz", layout_gauss); + 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); } void WgPipelines::releaseGraphicHandles(WgContext& context) { // pipeline effects + releaseComputePipeline(fill_effect); releaseComputePipeline(dropshadow); releaseComputePipeline(gaussian_vert); releaseComputePipeline(gaussian_horz); @@ -492,8 +493,8 @@ void WgPipelines::releaseGraphicHandles(WgContext& context) releaseRenderPipeline(evenodd); releaseRenderPipeline(nonzero); // layouts - releasePipelineLayout(layout_dropshadow); - releasePipelineLayout(layout_gaussian); + releasePipelineLayout(layout_effects); + releasePipelineLayout(layout_gauss); releasePipelineLayout(layout_blit); releasePipelineLayout(layout_scene_compose); releasePipelineLayout(layout_scene_blend); @@ -507,8 +508,7 @@ void WgPipelines::releaseGraphicHandles(WgContext& context) releasePipelineLayout(layout_depth); releasePipelineLayout(layout_stencil); // shaders - releaseShaderModule(shader_dropshadow); - releaseShaderModule(shader_gaussian); + releaseShaderModule(shader_effects); 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 215376b2..73139e41 100755 --- a/src/renderer/wg_engine/tvgWgPipelines.h +++ b/src/renderer/wg_engine/tvgWgPipelines.h @@ -47,8 +47,8 @@ private: // shader blit WGPUShaderModule shader_blit{}; // shader effects - WGPUShaderModule shader_gaussian; - WGPUShaderModule shader_dropshadow; + WGPUShaderModule shader_gauss; + WGPUShaderModule shader_effects; // layouts helpers WGPUPipelineLayout layout_stencil{}; @@ -68,8 +68,8 @@ private: // layouts blit WGPUPipelineLayout layout_blit{}; // layouts effects - WGPUPipelineLayout layout_gaussian{}; - WGPUPipelineLayout layout_dropshadow{}; + WGPUPipelineLayout layout_gauss{}; + WGPUPipelineLayout layout_effects{}; public: // pipelines stencil markup WGPURenderPipeline nonzero{}; @@ -101,6 +101,7 @@ public: WGPUComputePipeline gaussian_horz{}; WGPUComputePipeline gaussian_vert{}; WGPUComputePipeline dropshadow{}; + WGPUComputePipeline fill_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 44f762d5..265a99f8 100755 --- a/src/renderer/wg_engine/tvgWgRenderData.cpp +++ b/src/renderer/wg_engine/tvgWgRenderData.cpp @@ -580,133 +580,81 @@ void WgRenderDataViewportPool::release(WgContext& context) } //*********************************************************************** -// WgRenderDataGaussian +// WgRenderDataEffectParams //*********************************************************************** -void WgRenderDataGaussian::update(WgContext& context, RenderEffectGaussianBlur* gaussian, const Matrix& transform) +void WgRenderDataEffectParams::update(WgContext& context, const WgShaderTypeEffectParams& effectParams) +{ + if (context.allocateBufferUniform(bufferParams, &effectParams.params, sizeof(effectParams.params))) { + context.layouts.releaseBindGroup(bindGroupParams); + bindGroupParams = context.layouts.createBindGroupBuffer1Un(bufferParams); + } +} + + +void WgRenderDataEffectParams::update(WgContext& context, const 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); - } + WgShaderTypeEffectParams effectParams; + effectParams.update(gaussian, transform); + update(context, effectParams); level = int(WG_GAUSSIAN_MAX_LEVEL * ((gaussian->quality - 1) * 0.01f)) + 1; - extend = gaussianSettings.extend; + extend = effectParams.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(); -} - -//*********************************************************************** -// WgRenderDataDropShadow -//*********************************************************************** - -void WgRenderDataDropShadow::update(WgContext& context, RenderEffectDropShadow* dropShadow, const Matrix& transform) +void WgRenderDataEffectParams::update(WgContext& context, const 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); - } + WgShaderTypeEffectParams effectParams; + effectParams.update(dropShadow, transform); + update(context, effectParams); level = int(WG_GAUSSIAN_MAX_LEVEL * ((dropShadow->quality - 1) * 0.01f)) + 1; - extend = gaussianSettings.extend; - offset = dropShadowSettings.offset; + extend = effectParams.extend; + offset = effectParams.offset; } -void WgRenderDataDropShadow::release(WgContext& context) +void WgRenderDataEffectParams::update(WgContext& context, const RenderEffectFill* fill) { - context.releaseBuffer(bufferSettings); - context.releaseBuffer(bufferGaussian); - context.layouts.releaseBindGroup(bindGroupDropShadow); - context.layouts.releaseBindGroup(bindGroupGaussian); + assert(fill); + WgShaderTypeEffectParams effectParams; + effectParams.update(fill); + update(context, effectParams); +} + + +void WgRenderDataEffectParams::release(WgContext& context) +{ + context.releaseBuffer(bufferParams); + context.layouts.releaseBindGroup(bindGroupParams); } //*********************************************************************** -// WgRenderDataGaussianPool +// WgRenderDataColorsPool //*********************************************************************** -WgRenderDataDropShadow* WgRenderDataDropShadowPool::allocate(WgContext& context) +WgRenderDataEffectParams* WgRenderDataEffectParamsPool::allocate(WgContext& context) { - WgRenderDataDropShadow* renderData{}; + WgRenderDataEffectParams* renderData{}; if (mPool.count > 0) { renderData = mPool.last(); mPool.pop(); } else { - renderData = new WgRenderDataDropShadow(); + renderData = new WgRenderDataEffectParams(); mList.push(renderData); } return renderData; } -void WgRenderDataDropShadowPool::free(WgContext& context, WgRenderDataDropShadow* renderData) +void WgRenderDataEffectParamsPool::free(WgContext& context, WgRenderDataEffectParams* renderData) { if (renderData) mPool.push(renderData); } -void WgRenderDataDropShadowPool::release(WgContext& context) +void WgRenderDataEffectParamsPool::release(WgContext& context) { ARRAY_FOREACH(p, mList) { (*p)->release(context); diff --git a/src/renderer/wg_engine/tvgWgRenderData.h b/src/renderer/wg_engine/tvgWgRenderData.h index 8f32aa93..9a2ab81e 100755 --- a/src/renderer/wg_engine/tvgWgRenderData.h +++ b/src/renderer/wg_engine/tvgWgRenderData.h @@ -25,6 +25,7 @@ #include "tvgWgPipelines.h" #include "tvgWgGeometry.h" +#include "tvgWgShaderTypes.h" struct WgMeshData { WGPUBuffer bufferPosition{}; @@ -190,55 +191,34 @@ public: void release(WgContext& context); }; +// gaussian blur, drop shadow, fill, tint, tritone #define WG_GAUSSIAN_MAX_LEVEL 3 -struct WgRenderDataGaussian +struct WgRenderDataEffectParams { - 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); -}; - -struct WgRenderDataDropShadow -{ - WGPUBindGroup bindGroupGaussian{}; - WGPUBindGroup bindGroupDropShadow{}; - WGPUBuffer bufferGaussian{}; - WGPUBuffer bufferSettings{}; + WGPUBindGroup bindGroupParams{}; + WGPUBuffer bufferParams{}; uint32_t extend{}; uint32_t level{}; Point offset{}; - void update(WgContext& context, RenderEffectDropShadow* dropShadow, const Matrix& transform); + void update(WgContext& context, const WgShaderTypeEffectParams& effectParams); + 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 release(WgContext& context); }; -class WgRenderDataDropShadowPool { +// effect params pool +class WgRenderDataEffectParamsPool { private: - // pool contains all created but unused render data for drop shadow - Array mPool; - // list contains all created render data for drop shadow + // pool contains all created but unused render data for params + Array mPool; + // list contains all created render data for params // to ensure that all created instances will be released - Array mList; + Array mList; public: - WgRenderDataDropShadow* allocate(WgContext& context); - void free(WgContext& context, WgRenderDataDropShadow* renderData); + WgRenderDataEffectParams* allocate(WgContext& context); + void free(WgContext& context, WgRenderDataEffectParams* renderData); void release(WgContext& context); }; diff --git a/src/renderer/wg_engine/tvgWgRenderer.cpp b/src/renderer/wg_engine/tvgWgRenderer.cpp index 34e12580..72d78405 100755 --- a/src/renderer/wg_engine/tvgWgRenderer.cpp +++ b/src/renderer/wg_engine/tvgWgRenderer.cpp @@ -52,9 +52,8 @@ void WgRenderer::release() // clear render data paint pools mRenderDataShapePool.release(mContext); mRenderDataPicturePool.release(mContext); - mRenderDataGaussianPool.release(mContext); - mRenderDataDropShadowPool.release(mContext); mRenderDataViewportPool.release(mContext); + mRenderDataEffectParamsPool.release(mContext); WgMeshDataPool::gMeshDataPool->release(mContext); // clear render storage pool @@ -524,24 +523,35 @@ void WgRenderer::prepare(RenderEffect* effect, const Matrix& transform) { // 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; + auto renderEffect = (RenderEffectGaussianBlur*)effect; + auto renderData = (WgRenderDataEffectParams*)renderEffect->rd; + if (!renderData) { + renderData = mRenderDataEffectParamsPool.allocate(mContext); + renderEffect->rd = renderData; } - renderDataGaussian->update(mContext, gaussianBlur, transform); + renderData->update(mContext, renderEffect, 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; + auto renderEffect = (RenderEffectDropShadow*)effect; + auto renderData = (WgRenderDataEffectParams*)renderEffect->rd; + if (!renderData) { + renderData = mRenderDataEffectParamsPool.allocate(mContext); + renderEffect->rd = renderData; } - renderDataDropShadow->update(mContext, dropShadow, transform); + renderData->update(mContext, renderEffect, transform); + effect->valid = true; + } + // prepare fill + if (effect->type == SceneEffect::Fill) { + auto renderEffect = (RenderEffectFill*)effect; + auto renderData = (WgRenderDataEffectParams*)renderEffect->rd; + if (!renderData) { + renderData = mRenderDataEffectParamsPool.allocate(mContext); + renderEffect->rd = renderData; + } + renderData->update(mContext, renderEffect); effect->valid = true; } } @@ -552,25 +562,25 @@ bool WgRenderer::region(RenderEffect* effect) // update gaussian blur region if (effect->type == SceneEffect::GaussianBlur) { auto gaussian = (RenderEffectGaussianBlur*)effect; - auto renderDataGaussian = (WgRenderDataGaussian*)gaussian->rd; + auto renderData = (WgRenderDataEffectParams*)gaussian->rd; if (gaussian->direction != 2) { - gaussian->extend.x = -renderDataGaussian->extend; - gaussian->extend.w = +renderDataGaussian->extend * 2; + gaussian->extend.x = -renderData->extend; + gaussian->extend.w = +renderData->extend * 2; } if (gaussian->direction != 1) { - gaussian->extend.y = -renderDataGaussian->extend; - gaussian->extend.h = +renderDataGaussian->extend * 2; + gaussian->extend.y = -renderData->extend; + gaussian->extend.h = +renderData->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; + auto renderData = (WgRenderDataEffectParams*)dropShadow->rd; + dropShadow->extend.x = -(renderData->extend + std::abs(renderData->offset.x)); + dropShadow->extend.w = +(renderData->extend + std::abs(renderData->offset.x)) * 2; + dropShadow->extend.y = -(renderData->extend + std::abs(renderData->offset.y)); + dropShadow->extend.h = +(renderData->extend + std::abs(renderData->offset.y)) * 2; return true; } return false; @@ -587,6 +597,7 @@ bool WgRenderer::render(RenderCompositor* cmp, const RenderEffect* effect, TVG_U 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); + case SceneEffect::Fill: return mCompositor.fillEffect(mContext, dst, (RenderEffectFill*)effect, comp); default: return false; } return false; @@ -595,11 +606,8 @@ bool WgRenderer::render(RenderCompositor* cmp, const RenderEffect* effect, TVG_U void WgRenderer::dispose(RenderEffect* effect) { - 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; - } + auto renderData = (WgRenderDataEffectParams*)effect->rd; + mRenderDataEffectParamsPool.free(mContext, renderData); effect->rd = nullptr; }; diff --git a/src/renderer/wg_engine/tvgWgRenderer.h b/src/renderer/wg_engine/tvgWgRenderer.h index ef7c3d7d..01f1e17f 100755 --- a/src/renderer/wg_engine/tvgWgRenderer.h +++ b/src/renderer/wg_engine/tvgWgRenderer.h @@ -85,9 +85,8 @@ private: // render data paint pools WgRenderDataShapePool mRenderDataShapePool; WgRenderDataPicturePool mRenderDataPicturePool; - WgRenderDataGaussianPool mRenderDataGaussianPool; - WgRenderDataDropShadowPool mRenderDataDropShadowPool; WgRenderDataViewportPool mRenderDataViewportPool; + WgRenderDataEffectParamsPool mRenderDataEffectParamsPool; // rendering context WgContext mContext; diff --git a/src/renderer/wg_engine/tvgWgShaderSrc.cpp b/src/renderer/wg_engine/tvgWgShaderSrc.cpp index d6f6246d..6f8abdd3 100755 --- a/src/renderer/wg_engine/tvgWgShaderSrc.cpp +++ b/src/renderer/wg_engine/tvgWgShaderSrc.cpp @@ -711,7 +711,7 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f { const char* cShaderSrc_GaussianBlur = 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(2) @binding(0) var settings: array; @group(3) @binding(0) var viewport: vec4f; const N: u32 = 128; @@ -728,9 +728,9 @@ fn gaussian(x: f32, sigma: f32) -> f32 { fn cs_main_horz(@builtin(global_invocation_id) gid: vec3u, @builtin(local_invocation_id) lid: vec3u) { // settings decode - let sigma = settings.x; - let scale = settings.y; - let size = i32(settings.z); + let sigma = settings[0].x; + let scale = settings[0].y; + let size = i32(settings[0].z); // viewport decode let xmin = i32(viewport.x); @@ -771,9 +771,9 @@ fn cs_main_horz(@builtin(global_invocation_id) gid: vec3u, fn cs_main_vert(@builtin(global_invocation_id) gid: vec3u, @builtin(local_invocation_id) lid: vec3u) { // settings decode - let sigma = settings.x; - let scale = settings.y; - let size = i32(settings.z); + let sigma = settings[0].x; + let scale = settings[0].y; + let size = i32(settings[0].z); // viewport decode let xmin = i32(viewport.x); @@ -811,19 +811,19 @@ fn cs_main_vert(@builtin(global_invocation_id) gid: vec3u, } )"; -const char* cShaderSrc_DropShadow = R"( +const char* cShaderSrc_Effects = 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(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) { +fn cs_main_drop_shadow(@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); + let voff = vec2i(settings[2].xy); // tex coord let uid = gid.xy + vmin; @@ -831,8 +831,23 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3u) { let orig = textureLoad(imageSrc, uid); let blur = textureLoad(imageSdw, oid); - let shad = settings[0] * blur.a; + let shad = settings[1] * blur.a; let color = orig + shad * (1.0 - orig.a); textureStore(imageTrg, uid.xy, color); } + +@compute @workgroup_size(128, 1) +fn cs_main_fill(@builtin(global_invocation_id) gid: vec3u) { + // decode viewport and settings + let vmin = vec2u(viewport.xy); + let vmax = vec2u(viewport.zw); + + // tex coord + let uid = gid.xy + vmin; + + let orig = textureLoad(imageSrc, uid); + let fill = settings[0]; + let color = fill * orig.a * fill.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 126f5459..cbf0c86b 100755 --- a/src/renderer/wg_engine/tvgWgShaderSrc.h +++ b/src/renderer/wg_engine/tvgWgShaderSrc.h @@ -46,6 +46,6 @@ extern const char* cShaderSrc_Blit; // compute shader sources: effects extern const char* cShaderSrc_GaussianBlur; -extern const char* cShaderSrc_DropShadow; +extern const char* cShaderSrc_Effects; #endif // _TVG_WG_SHEDER_SRC_H_ diff --git a/src/renderer/wg_engine/tvgWgShaderTypes.cpp b/src/renderer/wg_engine/tvgWgShaderTypes.cpp index 8f0bc2f5..3e42c0bd 100755 --- a/src/renderer/wg_engine/tvgWgShaderTypes.cpp +++ b/src/renderer/wg_engine/tvgWgShaderTypes.cpp @@ -197,52 +197,53 @@ void WgShaderTypeGradient::updateTexData(const Fill::ColorStop* stops, uint32_t } //************************************************************************ -// WgShaderTypeGaussianBlur +// WgShaderTypeEffectParams //************************************************************************ -void WgShaderTypeGaussianBlur::update(float sigma, const Matrix& transform) -{ - 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; - settings[3] = 0.0f; // unused - extend = settings[2] * 2; // kernel -} - - -void WgShaderTypeGaussianBlur::update(const RenderEffectGaussianBlur* gaussian, const Matrix& transform) +void WgShaderTypeEffectParams::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 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 + params[0] = sigma; + params[1] = std::min(WG_GAUSSIAN_KERNEL_SIZE_MAX / kernel, scale); + params[2] = kernel; + params[3] = 0.0f; + extend = params[2] * 2; // kernel +} + + +void WgShaderTypeEffectParams::update(const RenderEffectDropShadow* dropShadow, const Matrix& transform) +{ + assert(dropShadow); const float radian = tvg::deg2rad(90.0f - dropShadow->angle); + const float sigma = dropShadow->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 offset = {0, 0}; - if (dropShadow->distance > 0.0f) offset = { + 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 + params[0] = sigma; + params[1] = std::min(WG_GAUSSIAN_KERNEL_SIZE_MAX / kernel, scale); + params[2] = kernel; + params[3] = 0.0f; + params[4] = dropShadow->color[0] / 255.0f; // red + params[5] = dropShadow->color[1] / 255.0f; // green + params[6] = dropShadow->color[2] / 255.0f; // blue + params[7] = dropShadow->color[3] / 255.0f; // alpha + params[8] = offset.x; + params[9] = offset.y; + extend = params[2] * 2; // kernel +} + + +void WgShaderTypeEffectParams::update(const RenderEffectFill* fill) +{ + params[0] = fill->color[0] / 255.0f; + params[1] = fill->color[1] / 255.0f; + params[2] = fill->color[2] / 255.0f; + params[3] = fill->color[3] / 255.0f; +} diff --git a/src/renderer/wg_engine/tvgWgShaderTypes.h b/src/renderer/wg_engine/tvgWgShaderTypes.h index 0bb3547c..e5aa969a 100755 --- a/src/renderer/wg_engine/tvgWgShaderTypes.h +++ b/src/renderer/wg_engine/tvgWgShaderTypes.h @@ -67,25 +67,21 @@ struct WgShaderTypeGradient void updateTexData(const Fill::ColorStop* stops, uint32_t stopCnt); }; -// gaussian settings: sigma, scale, extend +// gaussian params: sigma, scale, extend #define WG_GAUSSIAN_KERNEL_SIZE_MAX (128.0f) -struct WgShaderTypeGaussianBlur +// gaussian blur, drop shadow, fill, tint, tritone +struct WgShaderTypeEffectParams { - float settings[4]{}; // [0]: sigma, [1]: scale, [2]: kernel size, [3]: unused - uint32_t extend{}; + // 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 + float params[4+4+4]{}; // settings: array; + uint32_t extend{}; // gaussian blur extend + Point offset{}; // drop shadow offset - 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); + void update(const RenderEffectFill* fill); }; #endif // _TVG_WG_SHADER_TYPES_H_