diff --git a/src/renderer/wg_engine/tvgWgCommon.cpp b/src/renderer/wg_engine/tvgWgCommon.cpp index 79cea21f..bb0ebfc2 100644 --- a/src/renderer/wg_engine/tvgWgCommon.cpp +++ b/src/renderer/wg_engine/tvgWgCommon.cpp @@ -596,36 +596,21 @@ void WgRenderPipeline::set(WGPURenderPassEncoder renderPassEncoder) WGPUBlendState WgRenderPipeline::makeBlendState(WgPipelineBlendType blendType) { WGPUBlendState blendState{}; - // src - if (blendType == WgPipelineBlendType::Src) { + // srcover + if (blendType == WgPipelineBlendType::SrcOver) { blendState.color.operation = WGPUBlendOperation_Add; blendState.color.srcFactor = WGPUBlendFactor_One; blendState.color.dstFactor = WGPUBlendFactor_Zero; - } else // normal - if (blendType == WgPipelineBlendType::Normal) { + // normal + } else if (blendType == WgPipelineBlendType::Normal) { blendState.color.operation = WGPUBlendOperation_Add; blendState.color.srcFactor = WGPUBlendFactor_One; blendState.color.dstFactor = WGPUBlendFactor_OneMinusSrcAlpha; - } else // add - if (blendType == WgPipelineBlendType::Add) { + // custom + } else if (blendType == WgPipelineBlendType::Custom) { blendState.color.operation = WGPUBlendOperation_Add; blendState.color.srcFactor = WGPUBlendFactor_One; - blendState.color.dstFactor = WGPUBlendFactor_One; - } else // mult - if (blendType == WgPipelineBlendType::Mult) { - blendState.color.operation = WGPUBlendOperation_Add; - blendState.color.srcFactor = WGPUBlendFactor_Dst; blendState.color.dstFactor = WGPUBlendFactor_Zero; - } else // min - if (blendType == WgPipelineBlendType::Min) { - blendState.color.operation = WGPUBlendOperation_Min; - blendState.color.srcFactor = WGPUBlendFactor_One; - blendState.color.dstFactor = WGPUBlendFactor_One; - } else // max - if (blendType == WgPipelineBlendType::Max) { - blendState.color.operation = WGPUBlendOperation_Max; - blendState.color.srcFactor = WGPUBlendFactor_One; - blendState.color.dstFactor = WGPUBlendFactor_One; } blendState.alpha = blendState.color; return blendState; diff --git a/src/renderer/wg_engine/tvgWgCommon.h b/src/renderer/wg_engine/tvgWgCommon.h index 9edac4fa..d0cdef3a 100644 --- a/src/renderer/wg_engine/tvgWgCommon.h +++ b/src/renderer/wg_engine/tvgWgCommon.h @@ -36,12 +36,9 @@ struct WgCompositor: public Compositor { }; enum class WgPipelineBlendType { - Src = 0, // S - Normal, // (Sa * S) + (255 - Sa) * D - Add, // (S + D) - Mult, // (S * D) - Min, // min(S, D) - Max // max(S, D) + SrcOver = 0, + Normal, + Custom }; struct WgPipelines; diff --git a/src/renderer/wg_engine/tvgWgPipelines.cpp b/src/renderer/wg_engine/tvgWgPipelines.cpp index bb183727..f38af79e 100644 --- a/src/renderer/wg_engine/tvgWgPipelines.cpp +++ b/src/renderer/wg_engine/tvgWgPipelines.cpp @@ -55,7 +55,7 @@ void WgPipelineFillShapeWinding::initialize(WGPUDevice device) auto pipelineLabel = "The render pipeline fill shape winding"; // allocate all pipeline handles - allocate(device, WgPipelineBlendType::Add, WGPUColorWriteMask_None, + allocate(device, WgPipelineBlendType::SrcOver, WGPUColorWriteMask_None, vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts), bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts), stencilFunctionFront, stencilOperationFront, stencilFunctionBack, stencilOperationBack, @@ -89,7 +89,7 @@ void WgPipelineFillShapeEvenOdd::initialize(WGPUDevice device) auto pipelineLabel = "The render pipeline fill shape Even Odd"; // allocate all pipeline handles - allocate(device, WgPipelineBlendType::Add, WGPUColorWriteMask_None, + allocate(device, WgPipelineBlendType::SrcOver, WGPUColorWriteMask_None, vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts), bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts), stencilFunctionFront, stencilOperationFront, stencilFunctionBack, stencilOperationBack, @@ -121,7 +121,7 @@ void WgPipelineFillStroke::initialize(WGPUDevice device) auto pipelineLabel = "The render pipeline fill stroke"; // allocate all pipeline handles - allocate(device, WgPipelineBlendType::Add, WGPUColorWriteMask_None, + allocate(device, WgPipelineBlendType::SrcOver, WGPUColorWriteMask_None, vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts), bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts), stencilFunction, stencilOperation, stencilFunction, stencilOperation, @@ -285,7 +285,7 @@ void WgPipelineClear::initialize(WGPUDevice device) } -void WgPipelineBlend::initialize(WGPUDevice device) +void WgPipelineBlend::initialize(WGPUDevice device, const char *shaderSource) { // bind groups and layouts WGPUBindGroupLayout bindGroupLayouts[] = { @@ -296,7 +296,6 @@ void WgPipelineBlend::initialize(WGPUDevice device) }; // sheder source and labels - auto shaderSource = cShaderSource_PipelineComputeBlend; auto shaderLabel = "The compute shader blend"; auto pipelineLabel = "The compute pipeline blend"; @@ -380,7 +379,7 @@ void WgPipelines::initialize(WgContext& context) fillShapeWinding.initialize(context.device); fillShapeEvenOdd.initialize(context.device); fillStroke.initialize(context.device); - for (uint8_t type = (uint8_t)WgPipelineBlendType::Src; type <= (uint8_t)WgPipelineBlendType::Max; type++) { + for (uint8_t type = (uint8_t)WgPipelineBlendType::SrcOver; type <= (uint8_t)WgPipelineBlendType::Custom; type++) { solid[type].initialize(context.device, (WgPipelineBlendType)type); linear[type].initialize(context.device, (WgPipelineBlendType)type); radial[type].initialize(context.device, (WgPipelineBlendType)type); @@ -388,7 +387,9 @@ void WgPipelines::initialize(WgContext& context) } // compute pipelines computeClear.initialize(context.device); - computeBlend.initialize(context.device); + computeBlendSolid.initialize(context.device, cShaderSource_PipelineComputeBlendSolid); + computeBlendGradient.initialize(context.device, cShaderSource_PipelineComputeBlendGradient); + computeBlendImage.initialize(context.device, cShaderSource_PipelineComputeBlendImage); computeCompose.initialize(context.device); computeComposeBlend.initialize(context.device); computeAntiAliasing.initialize(context.device); @@ -415,10 +416,12 @@ void WgPipelines::release() computeAntiAliasing.release(); computeComposeBlend.release(); computeCompose.release(); - computeBlend.release(); + computeBlendImage.release(); + computeBlendGradient.release(); + computeBlendSolid.release(); computeClear.release(); // fill pipelines - for (uint8_t type = (uint8_t)WgPipelineBlendType::Src; type <= (uint8_t)WgPipelineBlendType::Max; type++) { + for (uint8_t type = (uint8_t)WgPipelineBlendType::SrcOver; type <= (uint8_t)WgPipelineBlendType::Custom; type++) { image[type].release(); radial[type].release(); linear[type].release(); @@ -435,10 +438,6 @@ bool WgPipelines::isBlendMethodSupportsHW(BlendMethod blendMethod) switch (blendMethod) { case BlendMethod::SrcOver: case BlendMethod::Normal: - case BlendMethod::Add: - case BlendMethod::Multiply: - case BlendMethod::Darken: - case BlendMethod::Lighten: return true; default: return false; }; @@ -448,12 +447,8 @@ bool WgPipelines::isBlendMethodSupportsHW(BlendMethod blendMethod) WgPipelineBlendType WgPipelines::blendMethodToBlendType(BlendMethod blendMethod) { switch (blendMethod) { - case BlendMethod::SrcOver: return WgPipelineBlendType::Src; + case BlendMethod::SrcOver: return WgPipelineBlendType::SrcOver; case BlendMethod::Normal: return WgPipelineBlendType::Normal; - case BlendMethod::Add: return WgPipelineBlendType::Add; - case BlendMethod::Multiply: return WgPipelineBlendType::Mult; - case BlendMethod::Darken: return WgPipelineBlendType::Min; - case BlendMethod::Lighten: return WgPipelineBlendType::Max; - default: return WgPipelineBlendType::Src; + default: return WgPipelineBlendType::Custom; }; } diff --git a/src/renderer/wg_engine/tvgWgPipelines.h b/src/renderer/wg_engine/tvgWgPipelines.h index d653f1fc..8b815af5 100644 --- a/src/renderer/wg_engine/tvgWgPipelines.h +++ b/src/renderer/wg_engine/tvgWgPipelines.h @@ -103,7 +103,7 @@ struct WgPipelineRadial: public WgRenderPipeline struct WgPipelineImage: public WgRenderPipeline { - void initialize(WGPUDevice device) override {} + void initialize(WGPUDevice device) override { assert(false); }; void initialize(WGPUDevice device, WgPipelineBlendType blendType); void use(WGPURenderPassEncoder encoder, WgBindGroupCanvas& groupCanvas, WgBindGroupPaint& groupPaint, WgBindGroupPicture& groupPicture) { @@ -131,7 +131,8 @@ struct WgPipelineClear: public WgComputePipeline struct WgPipelineBlend: public WgComputePipeline { - void initialize(WGPUDevice device) override; + void initialize(WGPUDevice device) override { assert(false); }; + void initialize(WGPUDevice device, const char *shaderSource); void use(WGPUComputePassEncoder encoder, WgBindGroupTextureStorageRgba& groupTexSrc, WgBindGroupTextureStorageRgba& groupTexDst, WgBindGroupBlendMethod& blendMethod, WgBindGroupOpacity& groupOpacity) { set(encoder); @@ -193,13 +194,15 @@ struct WgPipelines WgPipelineFillShapeEvenOdd fillShapeEvenOdd; WgPipelineFillStroke fillStroke; // fill pipelines - WgPipelineSolid solid[6]; - WgPipelineLinear linear[6]; - WgPipelineRadial radial[6]; - WgPipelineImage image[6]; + WgPipelineSolid solid[3]; + WgPipelineLinear linear[3]; + WgPipelineRadial radial[3]; + WgPipelineImage image[3]; // compute pipelines WgPipelineClear computeClear; - WgPipelineBlend computeBlend; + WgPipelineBlend computeBlendSolid; + WgPipelineBlend computeBlendGradient; + WgPipelineBlend computeBlendImage; WgPipelineCompose computeCompose; WgPipelineComposeBlend computeComposeBlend; WgPipelineAntiAliasing computeAntiAliasing; diff --git a/src/renderer/wg_engine/tvgWgRenderTarget.cpp b/src/renderer/wg_engine/tvgWgRenderTarget.cpp index 2f890681..a6a0a245 100644 --- a/src/renderer/wg_engine/tvgWgRenderTarget.cpp +++ b/src/renderer/wg_engine/tvgWgRenderTarget.cpp @@ -173,15 +173,15 @@ void WgRenderStorage::clear(WGPUCommandEncoder commandEncoder) } -void WgRenderStorage::blend(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc, WgBindGroupBlendMethod* blendMethod, WgBindGroupOpacity* opacity) -{ +void WgRenderStorage::blend(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc, WgPipelineBlend* pipeline, WgBindGroupBlendMethod* blendMethod, WgBindGroupOpacity* opacity) { assert(commandEncoder); assert(targetSrc); + assert(pipeline); WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder); - mPipelines->computeBlend.use(computePassEncoder, targetSrc->bindGroupTexStorageRgba, bindGroupTexStorageRgba, *blendMethod, *opacity); + pipeline->use(computePassEncoder, targetSrc->bindGroupTexStorageRgba, bindGroupTexStorageRgba, *blendMethod, *opacity); dispatchWorkgroups(computePassEncoder); endComputePass(computePassEncoder); -}; +} void WgRenderStorage::compose(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetMsk, WgBindGroupCompositeMethod* composeMethod, WgBindGroupOpacity* opacity) diff --git a/src/renderer/wg_engine/tvgWgRenderTarget.h b/src/renderer/wg_engine/tvgWgRenderTarget.h index 6e1108cb..527505d6 100644 --- a/src/renderer/wg_engine/tvgWgRenderTarget.h +++ b/src/renderer/wg_engine/tvgWgRenderTarget.h @@ -53,7 +53,7 @@ public: void renderPicture(WgContext& context, WgRenderDataPicture* renderData, WgPipelineBlendType blendType); void clear(WGPUCommandEncoder commandEncoder); - void blend(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc, WgBindGroupBlendMethod* blendMethod, WgBindGroupOpacity* opacity); + void blend(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc, WgPipelineBlend* pipeline, WgBindGroupBlendMethod* blendMethod, WgBindGroupOpacity* opacity); void compose(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetMsk, WgBindGroupCompositeMethod* composeMethod, WgBindGroupOpacity* opacity); void composeBlend( WgContext& context, diff --git a/src/renderer/wg_engine/tvgWgRenderer.cpp b/src/renderer/wg_engine/tvgWgRenderer.cpp index c73ee834..00eddf94 100644 --- a/src/renderer/wg_engine/tvgWgRenderer.cpp +++ b/src/renderer/wg_engine/tvgWgRenderer.cpp @@ -63,7 +63,7 @@ void WgRenderer::release() mOpacityPool.release(mContext); mRenderStorageRoot.release(mContext); mRenderStorageScreen.release(mContext); - mRenderTarget.release(mContext); + mRenderStorageInterm.release(mContext); mPipelines.release(); } @@ -152,24 +152,29 @@ bool WgRenderer::preRender() bool WgRenderer::renderShape(RenderData data) { // get current render storage - if (((WgRenderDataShape*)data)->opacity == 0) return 0; + WgRenderDataShape *dataShape = (WgRenderDataShape*)data; + if (dataShape->opacity == 0) return 0; WgPipelineBlendType blendType = WgPipelines::blendMethodToBlendType(mBlendMethod); WgRenderStorage* renderStorage = mRenderStorageStack.last(); assert(renderStorage); // use hardware blend if (WgPipelines::isBlendMethodSupportsHW(mBlendMethod)) - renderStorage->renderShape(mContext, (WgRenderDataShape *)data, blendType); - else { // use custom blend + renderStorage->renderShape(mContext, dataShape, blendType); + // use custom blend + else { // terminate current render pass renderStorage->endRenderPass(); // render image to render target - mRenderTarget.beginRenderPass(mCommandEncoder, true); - mRenderTarget.renderShape(mContext, (WgRenderDataShape *)data, blendType); - mRenderTarget.endRenderPass(); + mRenderStorageInterm.beginRenderPass(mCommandEncoder, true); + mRenderStorageInterm.renderShape(mContext, dataShape, blendType); + mRenderStorageInterm.endRenderPass(); // blend shape with current render storage WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod); WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, 255); - renderStorage->blend(mCommandEncoder, &mRenderTarget, blendMethod, opacity); + WgPipelineBlend* pipeline = &mContext.pipelines->computeBlendSolid; + if (dataShape->renderSettingsShape.fillType != WgRenderSettingsType::Solid) + pipeline = &mContext.pipelines->computeBlendGradient; + renderStorage->blend(mCommandEncoder, &mRenderStorageInterm, pipeline, blendMethod, opacity); // restore current render pass renderStorage->beginRenderPass(mCommandEncoder, false); } @@ -186,17 +191,19 @@ bool WgRenderer::renderImage(RenderData data) // use hardware blend if (WgPipelines::isBlendMethodSupportsHW(mBlendMethod)) renderStorage->renderPicture(mContext, (WgRenderDataPicture *)data, blendType); - else { // use custom blend + // use custom blend + else { // terminate current render pass renderStorage->endRenderPass(); // render image to render target - mRenderTarget.beginRenderPass(mCommandEncoder, true); - mRenderTarget.renderPicture(mContext, (WgRenderDataPicture *)data, blendType); - mRenderTarget.endRenderPass(); + mRenderStorageInterm.beginRenderPass(mCommandEncoder, true); + mRenderStorageInterm.renderPicture(mContext, (WgRenderDataPicture *)data, blendType); + mRenderStorageInterm.endRenderPass(); // blend shape with current render storage WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod); WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, 255); - renderStorage->blend(mCommandEncoder, &mRenderTarget, blendMethod, opacity); + WgPipelineBlend* pipeline = &mContext.pipelines->computeBlendImage; + renderStorage->blend(mCommandEncoder, &mRenderStorageInterm, pipeline, blendMethod, opacity); // restore current render pass renderStorage->beginRenderPass(mCommandEncoder, false); } @@ -293,7 +300,6 @@ bool WgRenderer::sync() mContext.executeCommandEncoder(commandEncoder); wgpuCommandEncoderRelease(commandEncoder); - wgpuSurfacePresent(mContext.surface); return true; } @@ -323,10 +329,9 @@ bool WgRenderer::target(WGPUInstance instance, WGPUSurface surface, uint32_t w, wgpuSurfaceConfigure(mContext.surface, &surfaceConfiguration); initialize(); - mRenderTarget.initialize(mContext, w, h, WG_SSAA_SAMPLES); + mRenderStorageInterm.initialize(mContext, w, h, WG_SSAA_SAMPLES); mRenderStorageRoot.initialize(mContext, w, h, WG_SSAA_SAMPLES); mRenderStorageScreen.initialize(mContext, w, h, 1, WGPUTextureFormat_BGRA8Unorm); - return true; } @@ -366,11 +371,11 @@ bool WgRenderer::endComposite(TVG_UNUSED Compositor* cmp) // get two last render targets WgRenderStorage* renderStorageSrc = mRenderStorageStack.last(); mRenderStorageStack.pop(); - + // blent scene to current render storage WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, comp->blendMethod); WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, comp->opacity); - mRenderStorageStack.last()->blend(mCommandEncoder, renderStorageSrc, blendMethod, opacity); + mRenderStorageStack.last()->blend(mCommandEncoder, renderStorageSrc, &mContext.pipelines->computeBlendImage, blendMethod, opacity); // back render targets to the pool mRenderStoragePool.free(mContext, renderStorageSrc); diff --git a/src/renderer/wg_engine/tvgWgRenderer.h b/src/renderer/wg_engine/tvgWgRenderer.h index 9757bbe7..e13fcdef 100644 --- a/src/renderer/wg_engine/tvgWgRenderer.h +++ b/src/renderer/wg_engine/tvgWgRenderer.h @@ -63,13 +63,19 @@ public: private: // render handles WGPUCommandEncoder mCommandEncoder{}; - WgRenderStorage mRenderTarget; + // intermidiate buffer to render + WgRenderStorage mRenderStorageInterm; + // root render storage WgRenderStorage mRenderStorageRoot; + // storage with data after antializing WgRenderStorage mRenderStorageScreen; + // pool to hold render tree storages WgRenderStoragePool mRenderStoragePool; + // opacity, blend methods and composite methods pool WgBindGroupOpacityPool mOpacityPool; WgBindGroupBlendMethodPool mBlendMethodPool; WgBindGroupCompositeMethodPool mCompositeMethodPool; + // render data shpes pool WgRenderDataShapePool mRenderDataShapePool; // render tree stacks @@ -80,6 +86,7 @@ private: WgContext mContext; WgPipelines mPipelines; Surface mTargetSurface; + // current blend method BlendMethod mBlendMethod{}; }; diff --git a/src/renderer/wg_engine/tvgWgShaderSrc.cpp b/src/renderer/wg_engine/tvgWgShaderSrc.cpp index 2ebec6fa..bea7bc77 100644 --- a/src/renderer/wg_engine/tvgWgShaderSrc.cpp +++ b/src/renderer/wg_engine/tvgWgShaderSrc.cpp @@ -23,11 +23,13 @@ #include "tvgWgShaderSrc.h" #include +#define WG_SHADER_SOURCE(...) #__VA_ARGS__ + //************************************************************************ // shader pipeline fill //************************************************************************ -const char* cShaderSource_PipelineFill = R"( +const char* cShaderSource_PipelineFill = WG_SHADER_SOURCE( // vertex input struct VertexInput { @location(0) position: vec2f @@ -54,13 +56,13 @@ fn vs_main(in: VertexInput) -> VertexOutput { fn fs_main(in: VertexOutput) -> void { // nothing to draw, just stencil value } -)"; +); //************************************************************************ // shader pipeline solid //************************************************************************ -const char* cShaderSource_PipelineSolid = R"( +const char* cShaderSource_PipelineSolid = WG_SHADER_SOURCE( // vertex input struct VertexInput { @location(0) position: vec2f @@ -80,10 +82,10 @@ struct VertexOutput { }; // uniforms -@group(0) @binding(0) var uViewMat : mat4x4f; -@group(1) @binding(0) var uModelMat : mat4x4f; +@group(0) @binding(0) var uViewMat : mat4x4f; +@group(1) @binding(0) var uModelMat : mat4x4f; @group(1) @binding(1) var uBlendSettings : BlendSettings; -@group(2) @binding(0) var uSolidColor : vec4f; +@group(2) @binding(0) var uSolidColor : vec4f; @vertex fn vs_main(in: VertexInput) -> VertexOutput { @@ -102,15 +104,15 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f { color = uSolidColor; let alpha: f32 = color.a * uBlendSettings.opacity; - return vec4f(color.rgb*alpha, alpha); + return vec4f(color.rgb * alpha, alpha); } -)"; +); //************************************************************************ // shader pipeline linear //************************************************************************ -const char* cShaderSource_PipelineLinear = R"( +const char* cShaderSource_PipelineLinear = WG_SHADER_SOURCE( // vertex input struct VertexInput { @location(0) position: vec2f @@ -171,7 +173,7 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f { let ba: vec2f = ed - st; // get interpolation factor - var t: f32 = abs(dot(pos - st, ba) / dot(ba, ba)); + var t: f32 = dot(pos - st, ba) / dot(ba, ba); // fill spread switch uLinearGradient.spread { @@ -205,15 +207,15 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f { } let alpha: f32 = color.a * uBlendSettings.opacity; - return vec4f(color.rgb*alpha, alpha); + return vec4f(color.rgb * alpha, alpha); } -)"; +); //************************************************************************ // shader pipeline radial //************************************************************************ -const char* cShaderSource_PipelineRadial = R"( +const char* cShaderSource_PipelineRadial = WG_SHADER_SOURCE( // vertex input struct VertexInput { @location(0) position: vec2f @@ -250,7 +252,7 @@ struct VertexOutput { // uniforms @group(0) @binding(0) var uViewMat : mat4x4f; @group(1) @binding(0) var uModelMat : mat4x4f; -@group(1) @binding(1) var uBlendSettings : BlendSettings; +@group(1) @binding(1) var uBlendSettings : BlendSettings; @group(2) @binding(0) var uRadialGradient : RadialGradient; @vertex @@ -302,15 +304,15 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f { } let alpha: f32 = color.a * uBlendSettings.opacity; - return vec4f(color.rgb*alpha, alpha); + return vec4f(color.rgb * alpha, alpha); } -)"; +); //************************************************************************ // cShaderSource_PipelineImage //************************************************************************ -const char* cShaderSource_PipelineImage = R"( +const char* cShaderSource_PipelineImage = WG_SHADER_SOURCE( // vertex input struct VertexInput { @location(0) position: vec2f, @@ -331,11 +333,11 @@ struct VertexOutput { @location(0) texCoord: vec2f }; -@group(0) @binding(0) var uViewMat : mat4x4f; -@group(1) @binding(0) var uModelMat : mat4x4f; +@group(0) @binding(0) var uViewMat : mat4x4f; +@group(1) @binding(0) var uModelMat : mat4x4f; @group(1) @binding(1) var uBlendSettings : BlendSettings; -@group(2) @binding(0) var uSampler : sampler; -@group(2) @binding(1) var uTextureView : texture_2d; +@group(2) @binding(0) var uSampler : sampler; +@group(2) @binding(1) var uTextureView : texture_2d; @vertex fn vs_main(in: VertexInput) -> VertexOutput { @@ -360,25 +362,13 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f { } return result * uBlendSettings.opacity; }; -)"; +); //************************************************************************ // cShaderSource_PipelineComputeBlend //************************************************************************ -// pipeline shader modules clear -const char* cShaderSource_PipelineComputeClear = R"( -@group(0) @binding(0) var imageDst : texture_storage_2d; - -@compute @workgroup_size(8, 8) -fn cs_main( @builtin(global_invocation_id) id: vec3u) { - textureStore(imageDst, id.xy, vec4f(0.0, 0.0, 0.0, 0.0)); -} -)"; - - -// pipeline shader modules blend -const char* cShaderSource_PipelineComputeBlend = R"( +const std::string strBlendShaderHeader = WG_SHADER_SOURCE( @group(0) @binding(0) var imageSrc : texture_storage_2d; @group(1) @binding(0) var imageDst : texture_storage_2d; @group(2) @binding(0) var blendMethod : u32; @@ -389,41 +379,123 @@ fn cs_main( @builtin(global_invocation_id) id: vec3u) { let texSize = textureDimensions(imageSrc); if ((id.x >= texSize.x) || (id.y >= texSize.y)) { return; }; - let colorSrc = textureLoad(imageSrc, id.xy); + let colorSrc= textureLoad(imageSrc, id.xy); + if (colorSrc.a == 0.0) { return; } let colorDst = textureLoad(imageDst, id.xy); - let So: f32 = opacity; - let Sa: f32 = colorSrc.a; - let Da: f32 = colorDst.a; - let S: vec4f = colorSrc; - let D: vec4f = colorDst; - let One: vec4f = vec4(1.0); + var One: vec3f = vec3(1.0); + var So: f32 = opacity; + var Sc: vec3f = colorSrc.rgb; + var Sa: f32 = colorSrc.a; + var Dc: vec3f = colorDst.rgb; + var Da: f32 = colorDst.a; + var Rc: vec3f = colorDst.rgb; + var Ra: f32 = 1.0; +); - var color: vec4f = vec4f(0.0, 0.0, 0.0, 0.0); +const std::string strBlendShaderPreConditionsGradient = WG_SHADER_SOURCE( + Sc = Sc + Dc.rgb * (1.0 - Sa); + Sa = Sa + Da * (1.0 - Sa); +); + +const std::string strBlendShaderPreConditionsImage = WG_SHADER_SOURCE( + Sc = Sc * So; + Sa = Sa * So; +); + +const std::string strBlendShaderBlendMethod = WG_SHADER_SOURCE( switch blendMethod { - /* Normal */ case 0u: { color = (S * So) + (1.0 - Sa * So) * D; } - /* Add */ case 1u: { color = (S + D); } - /* Screen */ case 2u: { color = (S + D) - (S * D); } - /* Multiply */ case 3u: { color = (S * D); } - /* Overlay */ case 4u: { color = (Sa * Da) - 2 * (Da - S) * (Sa - D); } - /* Difference */ case 5u: { color = abs(S - D); } - /* Exclusion */ case 6u: { color = S + D - (2 * S * D); } - /* SrcOver */ case 7u: { color = S; } - /* Darken */ case 8u: { color = min(S, D); } - /* Lighten */ case 9u: { color = max(S, D); } - /* ColorDodge */ case 10u: { color = D / (One - S); } - /* ColorBurn */ case 11u: { color = One - (One - D) / S; } - /* HardLight */ case 12u: { color = (Sa * Da) - 2.0 * (Da - S) * (Sa - D); } - /* SoftLight */ case 13u: { color = (One - 2 * S) * (D * D) + (2 * S * D); } - default: { color = (S * So) + (1.0 - Sa * So) * D; } + /* Normal */ case 0u: { + Rc = Sc + Dc * (1.0 - Sa); + Ra = Sa + Da * (1.0 - Sa); + } + /* Add */ case 1u: { Rc = Sc + Dc; } + /* Screen */ case 2u: { Rc = Sc + Dc - Sc * Dc; } + /* Multiply */ case 3u: { Rc = Sc * Dc; } + /* Overlay */ case 4u: { + Rc.r = select(1.0 - min(1.0, 2 * (1 - Sc.r) * (1 - Dc.r)), min(1.0, 2 * Sc.r * Dc.r), (Dc.r < 0.5)); + Rc.g = select(1.0 - min(1.0, 2 * (1 - Sc.g) * (1 - Dc.g)), min(1.0, 2 * Sc.g * Dc.g), (Dc.g < 0.5)); + Rc.b = select(1.0 - min(1.0, 2 * (1 - Sc.b) * (1 - Dc.b)), min(1.0, 2 * Sc.b * Dc.b), (Dc.b < 0.5)); + } + /* Difference */ case 5u: { Rc = abs(Dc - Sc); } + /* Exclusion */ case 6u: { Rc = min(One, Sc + Dc - min(One, 2 * Sc * Dc)); } + /* SrcOver */ case 7u: { Rc = Sc; Ra = Sa; } + /* Darken */ case 8u: { Rc = min(Sc, Dc); } + /* Lighten */ case 9u: { Rc = max(Sc, Dc); } + /* ColorDodge */ case 10u: { + Rc.r = select(Dc.r, (Dc.r * 255.0 / (255.0 - Sc.r * 255.0))/255.0, (1.0 - Sc.r > 0.0)); + Rc.g = select(Dc.g, (Dc.g * 255.0 / (255.0 - Sc.g * 255.0))/255.0, (1.0 - Sc.g > 0.0)); + Rc.b = select(Dc.b, (Dc.b * 255.0 / (255.0 - Sc.b * 255.0))/255.0, (1.0 - Sc.b > 0.0)); + } + /* ColorBurn */ case 11u: { + Rc.r = select(1.0 - Dc.r, (255.0 - (255.0 - Dc.r * 255.0) / (Sc.r * 255.0)) / 255.0, (Sc.r > 0.0)); + Rc.g = select(1.0 - Dc.g, (255.0 - (255.0 - Dc.g * 255.0) / (Sc.g * 255.0)) / 255.0, (Sc.g > 0.0)); + Rc.b = select(1.0 - Dc.b, (255.0 - (255.0 - Dc.b * 255.0) / (Sc.b * 255.0)) / 255.0, (Sc.b > 0.0)); + } + /* HardLight */ case 12u: { + Rc.r = select(1.0 - min(1.0, 2 * (1 - Sc.r) * (1 - Dc.r)), min(1.0, 2 * Sc.r * Dc.r), (Sc.r < 0.5)); + Rc.g = select(1.0 - min(1.0, 2 * (1 - Sc.g) * (1 - Dc.g)), min(1.0, 2 * Sc.g * Dc.g), (Sc.g < 0.5)); + Rc.b = select(1.0 - min(1.0, 2 * (1 - Sc.b) * (1 - Dc.b)), min(1.0, 2 * Sc.b * Dc.b), (Sc.b < 0.5)); + } + /* SoftLight */ case 13u: { Rc = min(One, (One - 2 * Sc) * Dc * Dc + 2.0 * Sc * Dc); } + default: { + Rc = Sc + Dc * (1.0 - Sa); + Ra = Sa + Da * (1.0 - Sa); + } } +); - textureStore(imageDst, id.xy, color); +const std::string strBlendShaderPostConditionsGradient = WG_SHADER_SOURCE( + // nothing +); + +const std::string strBlendShaderPostConditionsImage = WG_SHADER_SOURCE( + Rc = select(mix(Dc, Rc, Sa), Rc, blendMethod == 0); + Ra = select(mix(Da, Ra, Sa), Ra, blendMethod == 0); +); + +const std::string strBlendShaderFooter = WG_SHADER_SOURCE( + textureStore(imageDst, id.xy, vec4(Rc, Ra)); } -)"; +); + +// pipeline shader modules blend solid +const std::string strComputeBlendSolid = + strBlendShaderHeader + + strBlendShaderBlendMethod + + strBlendShaderFooter; +const char* cShaderSource_PipelineComputeBlendSolid = strComputeBlendSolid.c_str(); + +// pipeline shader modules blend gradient +const std::string strComputeBlendGradient = + strBlendShaderHeader + + strBlendShaderPreConditionsGradient + + strBlendShaderBlendMethod + + strBlendShaderPostConditionsGradient + + strBlendShaderFooter; +const char* cShaderSource_PipelineComputeBlendGradient = strComputeBlendGradient.c_str(); + +// pipeline shader modules blend image +const std::string strComputeBlendImage = + strBlendShaderHeader + + strBlendShaderPreConditionsImage + + strBlendShaderBlendMethod + + strBlendShaderPostConditionsImage + + strBlendShaderFooter; +const char* cShaderSource_PipelineComputeBlendImage = strComputeBlendImage.c_str(); + +// pipeline shader modules clear +const char* cShaderSource_PipelineComputeClear = WG_SHADER_SOURCE( +@group(0) @binding(0) var imageDst : texture_storage_2d; + +@compute @workgroup_size(8, 8) +fn cs_main( @builtin(global_invocation_id) id: vec3u) { + textureStore(imageDst, id.xy, vec4f(0.0, 0.0, 0.0, 0.0)); +} +); // pipeline shader modules compose -const char* cShaderSource_PipelineComputeCompose = R"( +const char* cShaderSource_PipelineComputeCompose = WG_SHADER_SOURCE( @group(0) @binding(0) var imageSrc : texture_storage_2d; @group(1) @binding(0) var imageMsk : texture_storage_2d; @group(2) @binding(0) var composeMethod : u32; @@ -455,10 +527,10 @@ fn cs_main( @builtin(global_invocation_id) id: vec3u) { textureStore(imageSrc, id.xy, vec4f(color, alpha * opacity)); } -)"; +); // pipeline shader modules compose blend -const char* cShaderSource_PipelineComputeComposeBlend = R"( +const char* cShaderSource_PipelineComputeComposeBlend = WG_SHADER_SOURCE( @group(0) @binding(0) var imageSrc : texture_storage_2d; @group(0) @binding(1) var imageMsk : texture_storage_2d; @group(0) @binding(2) var imageDst : texture_storage_2d; @@ -521,10 +593,10 @@ fn cs_main( @builtin(global_invocation_id) id: vec3u) { textureStore(imageDst, id.xy, color); } -)"; +); // pipeline shader modules anti-aliasing -const char* cShaderSource_PipelineComputeAntiAlias = R"( +const char* cShaderSource_PipelineComputeAntiAlias = WG_SHADER_SOURCE( @group(0) @binding(0) var imageSrc : texture_storage_2d; @group(1) @binding(0) var imageDst : texture_storage_2d; @@ -544,4 +616,4 @@ fn cs_main( @builtin(global_invocation_id) id: vec3u) { textureStore(imageDst, id.xy, color / f32(samples * samples)); } -)"; +); diff --git a/src/renderer/wg_engine/tvgWgShaderSrc.h b/src/renderer/wg_engine/tvgWgShaderSrc.h index 60d7c83b..abfc9d21 100644 --- a/src/renderer/wg_engine/tvgWgShaderSrc.h +++ b/src/renderer/wg_engine/tvgWgShaderSrc.h @@ -42,7 +42,9 @@ extern const char* cShaderSource_PipelineImage; // pipeline shader modules clear, compose and blend extern const char* cShaderSource_PipelineComputeClear; -extern const char* cShaderSource_PipelineComputeBlend; +extern const char* cShaderSource_PipelineComputeBlendSolid; +extern const char* cShaderSource_PipelineComputeBlendGradient; +extern const char* cShaderSource_PipelineComputeBlendImage; extern const char* cShaderSource_PipelineComputeCompose; extern const char* cShaderSource_PipelineComputeComposeBlend; extern const char* cShaderSource_PipelineComputeAntiAlias;