From 21286f679df98b4d38d21812ff87027a7ba1fae8 Mon Sep 17 00:00:00 2001 From: Sergii Liebodkin Date: Wed, 3 Jul 2024 07:47:08 +0300 Subject: [PATCH] wg_engine: fix scene opacity usage Changed blend equation to make it the same as for referenced sw renderer. --- src/renderer/wg_engine/tvgWgCommon.cpp | 6 +- src/renderer/wg_engine/tvgWgPipelines.cpp | 3 +- src/renderer/wg_engine/tvgWgPipelines.h | 3 +- src/renderer/wg_engine/tvgWgRenderTarget.cpp | 4 +- src/renderer/wg_engine/tvgWgRenderTarget.h | 2 +- src/renderer/wg_engine/tvgWgRenderer.cpp | 11 ++- src/renderer/wg_engine/tvgWgShaderSrc.cpp | 77 +++++++++++--------- 7 files changed, 58 insertions(+), 48 deletions(-) diff --git a/src/renderer/wg_engine/tvgWgCommon.cpp b/src/renderer/wg_engine/tvgWgCommon.cpp index f2588ff3..79cea21f 100644 --- a/src/renderer/wg_engine/tvgWgCommon.cpp +++ b/src/renderer/wg_engine/tvgWgCommon.cpp @@ -604,7 +604,7 @@ WGPUBlendState WgRenderPipeline::makeBlendState(WgPipelineBlendType blendType) } else // normal if (blendType == WgPipelineBlendType::Normal) { blendState.color.operation = WGPUBlendOperation_Add; - blendState.color.srcFactor = WGPUBlendFactor_SrcAlpha; + blendState.color.srcFactor = WGPUBlendFactor_One; blendState.color.dstFactor = WGPUBlendFactor_OneMinusSrcAlpha; } else // add if (blendType == WgPipelineBlendType::Add) { @@ -627,9 +627,7 @@ WGPUBlendState WgRenderPipeline::makeBlendState(WgPipelineBlendType blendType) blendState.color.srcFactor = WGPUBlendFactor_One; blendState.color.dstFactor = WGPUBlendFactor_One; } - blendState.alpha.operation = WGPUBlendOperation_Add; - blendState.alpha.srcFactor = WGPUBlendFactor_SrcAlpha; - blendState.alpha.dstFactor = WGPUBlendFactor_Zero; + blendState.alpha = blendState.color; return blendState; } diff --git a/src/renderer/wg_engine/tvgWgPipelines.cpp b/src/renderer/wg_engine/tvgWgPipelines.cpp index 2fcb9245..bb183727 100644 --- a/src/renderer/wg_engine/tvgWgPipelines.cpp +++ b/src/renderer/wg_engine/tvgWgPipelines.cpp @@ -291,7 +291,8 @@ void WgPipelineBlend::initialize(WGPUDevice device) WGPUBindGroupLayout bindGroupLayouts[] = { WgBindGroupTextureStorageRgba::getLayout(device), WgBindGroupTextureStorageRgba::getLayout(device), - WgBindGroupBlendMethod::getLayout(device) + WgBindGroupBlendMethod::getLayout(device), + WgBindGroupOpacity::getLayout(device) }; // sheder source and labels diff --git a/src/renderer/wg_engine/tvgWgPipelines.h b/src/renderer/wg_engine/tvgWgPipelines.h index 60e115cb..d653f1fc 100644 --- a/src/renderer/wg_engine/tvgWgPipelines.h +++ b/src/renderer/wg_engine/tvgWgPipelines.h @@ -132,12 +132,13 @@ struct WgPipelineClear: public WgComputePipeline struct WgPipelineBlend: public WgComputePipeline { void initialize(WGPUDevice device) override; - void use(WGPUComputePassEncoder encoder, WgBindGroupTextureStorageRgba& groupTexSrc, WgBindGroupTextureStorageRgba& groupTexDst, WgBindGroupBlendMethod& blendMethod) + void use(WGPUComputePassEncoder encoder, WgBindGroupTextureStorageRgba& groupTexSrc, WgBindGroupTextureStorageRgba& groupTexDst, WgBindGroupBlendMethod& blendMethod, WgBindGroupOpacity& groupOpacity) { set(encoder); groupTexSrc.set(encoder, 0); groupTexDst.set(encoder, 1); blendMethod.set(encoder, 2); + groupOpacity.set(encoder, 3); } }; diff --git a/src/renderer/wg_engine/tvgWgRenderTarget.cpp b/src/renderer/wg_engine/tvgWgRenderTarget.cpp index f39dae0b..2f890681 100644 --- a/src/renderer/wg_engine/tvgWgRenderTarget.cpp +++ b/src/renderer/wg_engine/tvgWgRenderTarget.cpp @@ -173,12 +173,12 @@ void WgRenderStorage::clear(WGPUCommandEncoder commandEncoder) } -void WgRenderStorage::blend(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc, WgBindGroupBlendMethod* blendMethod) +void WgRenderStorage::blend(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc, WgBindGroupBlendMethod* blendMethod, WgBindGroupOpacity* opacity) { assert(commandEncoder); assert(targetSrc); WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder); - mPipelines->computeBlend.use(computePassEncoder, targetSrc->bindGroupTexStorageRgba, bindGroupTexStorageRgba, *blendMethod); + mPipelines->computeBlend.use(computePassEncoder, targetSrc->bindGroupTexStorageRgba, bindGroupTexStorageRgba, *blendMethod, *opacity); dispatchWorkgroups(computePassEncoder); endComputePass(computePassEncoder); }; diff --git a/src/renderer/wg_engine/tvgWgRenderTarget.h b/src/renderer/wg_engine/tvgWgRenderTarget.h index 76e26629..6e1108cb 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); + void blend(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc, 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 2fa0ffa5..c73ee834 100644 --- a/src/renderer/wg_engine/tvgWgRenderer.cpp +++ b/src/renderer/wg_engine/tvgWgRenderer.cpp @@ -168,7 +168,8 @@ bool WgRenderer::renderShape(RenderData data) mRenderTarget.endRenderPass(); // blend shape with current render storage WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod); - renderStorage->blend(mCommandEncoder, &mRenderTarget, blendMethod); + WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, 255); + renderStorage->blend(mCommandEncoder, &mRenderTarget, blendMethod, opacity); // restore current render pass renderStorage->beginRenderPass(mCommandEncoder, false); } @@ -194,7 +195,8 @@ bool WgRenderer::renderImage(RenderData data) mRenderTarget.endRenderPass(); // blend shape with current render storage WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod); - renderStorage->blend(mCommandEncoder, &mRenderTarget, blendMethod); + WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, 255); + renderStorage->blend(mCommandEncoder, &mRenderTarget, blendMethod, opacity); // restore current render pass renderStorage->beginRenderPass(mCommandEncoder, false); } @@ -290,7 +292,7 @@ bool WgRenderer::sync() mContext.executeCommandEncoder(commandEncoder); wgpuCommandEncoderRelease(commandEncoder); - + wgpuSurfacePresent(mContext.surface); return true; } @@ -367,7 +369,8 @@ bool WgRenderer::endComposite(TVG_UNUSED Compositor* cmp) // blent scene to current render storage WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, comp->blendMethod); - mRenderStorageStack.last()->blend(mCommandEncoder, renderStorageSrc, blendMethod); + WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, comp->opacity); + mRenderStorageStack.last()->blend(mCommandEncoder, renderStorageSrc, blendMethod, opacity); // back render targets to the pool mRenderStoragePool.free(mContext, renderStorageSrc); diff --git a/src/renderer/wg_engine/tvgWgShaderSrc.cpp b/src/renderer/wg_engine/tvgWgShaderSrc.cpp index dcc72370..2ebec6fa 100644 --- a/src/renderer/wg_engine/tvgWgShaderSrc.cpp +++ b/src/renderer/wg_engine/tvgWgShaderSrc.cpp @@ -101,7 +101,8 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f { // get color color = uSolidColor; - return vec4f(color.rgb, color.a * uBlendSettings.opacity); + let alpha: f32 = color.a * uBlendSettings.opacity; + return vec4f(color.rgb*alpha, alpha); } )"; @@ -203,7 +204,8 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f { } } - return vec4f(color.rgb, color.a * uBlendSettings.opacity); + let alpha: f32 = color.a * uBlendSettings.opacity; + return vec4f(color.rgb*alpha, alpha); } )"; @@ -299,7 +301,8 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f { } } - return vec4f(color.rgb, color.a * uBlendSettings.opacity); + let alpha: f32 = color.a * uBlendSettings.opacity; + return vec4f(color.rgb*alpha, alpha); } )"; @@ -351,11 +354,11 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f { if (format == 1u) { /* FMT_ARGB8888 */ result = color.bgra; } else if (format == 2u) { /* FMT_ABGR8888S */ - result = vec4(color.rgb * color.a, color.a); + result = color.rgba; } else if (format == 3u) { /* FMT_ARGB8888S */ - result = vec4(color.bgr * color.a, color.a); + result = color.bgra; } - return vec4f(result.rgb, result.a * uBlendSettings.opacity); + return result * uBlendSettings.opacity; }; )"; @@ -379,25 +382,26 @@ const char* cShaderSource_PipelineComputeBlend = R"( @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; +@group(3) @binding(0) var opacity : f32; @compute @workgroup_size(8, 8) fn cs_main( @builtin(global_invocation_id) id: vec3u) { let texSize = textureDimensions(imageSrc); if ((id.x >= texSize.x) || (id.y >= texSize.y)) { return; }; - let srcColor = textureLoad(imageSrc, id.xy); - if (srcColor.a == 0.0) { return; }; - let dstColor = textureLoad(imageDst, id.xy); + let colorSrc = textureLoad(imageSrc, id.xy); + let colorDst = textureLoad(imageDst, id.xy); - let Sa: f32 = srcColor.a; - let Da: f32 = dstColor.a; - let S: vec3f = srcColor.xyz; - let D: vec3f = dstColor.xyz; - let One: vec3f = vec3(1.0); + 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 color: vec3f = vec3f(0.0, 0.0, 0.0); + var color: vec4f = vec4f(0.0, 0.0, 0.0, 0.0); switch blendMethod { - /* Normal */ case 0u: { color = (Sa * S) + (1.0 - Sa) * D; } + /* 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); } @@ -411,10 +415,10 @@ fn cs_main( @builtin(global_invocation_id) id: vec3u) { /* 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 = (Sa * S) + (1.0 - Sa) * D; } + default: { color = (S * So) + (1.0 - Sa * So) * D; } } - textureStore(imageDst, id.xy, vec4f(color, Sa)); + textureStore(imageDst, id.xy, color); } )"; @@ -472,30 +476,33 @@ fn cs_main( @builtin(global_invocation_id) id: vec3u) { if ((length(colorSrc) == 0.0) || (length(colorMsk) == 0.0)) { return; }; let colorDst = textureLoad(imageDst, id.xy); - var color: vec3f = colorSrc.xyz; + var result: vec3f = colorSrc.xyz; var alpha: f32 = colorMsk.a; let luma: f32 = dot(colorMsk.xyz, vec3f(0.299, 0.587, 0.114)); switch composeMethod { - /* None */ case 0u: { color = colorSrc.xyz; } + /* None */ case 0u: { result = colorSrc.xyz; } /* ClipPath */ case 1u: { if (colorMsk.a == 0) { alpha = 0.0; }; } - /* AlphaMask */ case 2u: { color = colorSrc.xyz; alpha = colorSrc.a * colorMsk.a; } - /* InvAlphaMask */ case 3u: { color = colorSrc.xyz; alpha = colorSrc.a * (1.0 - colorMsk.a); } - /* LumaMask */ case 4u: { color = colorSrc.xyz; alpha = colorSrc.a * luma; } - /* InvLumaMask */ case 5u: { color = colorSrc.xyz; alpha = colorSrc.a * (1.0 - luma); } - /* AddMask */ case 6u: { color = colorSrc.xyz * colorSrc.a + colorMsk.xyz * (1.0 - colorSrc.a); } - /* SubtractMask */ case 7u: { color = colorSrc.xyz * colorSrc.a - colorMsk.xyz * (1.0 - colorSrc.a); } - /* IntersectMask */ case 8u: { color = colorSrc.xyz * min(colorSrc.a, colorMsk.a); } - /* DifferenceMask */ case 9u: { color = abs(colorMsk.xyz - colorSrc.xyz * (1.0 - colorMsk.a)); } - default: { color = colorSrc.xyz; } + /* AlphaMask */ case 2u: { result = colorSrc.xyz; alpha = colorSrc.a * colorMsk.a; } + /* InvAlphaMask */ case 3u: { result = colorSrc.xyz; alpha = colorSrc.a * (1.0 - colorMsk.a); } + /* LumaMask */ case 4u: { result = colorSrc.xyz; alpha = colorSrc.a * luma; } + /* InvLumaMask */ case 5u: { result = colorSrc.xyz; alpha = colorSrc.a * (1.0 - luma); } + /* AddMask */ case 6u: { result = colorSrc.xyz * colorSrc.a + colorMsk.xyz * (1.0 - colorSrc.a); } + /* SubtractMask */ case 7u: { result = colorSrc.xyz * colorSrc.a - colorMsk.xyz * (1.0 - colorSrc.a); } + /* IntersectMask */ case 8u: { result = colorSrc.xyz * min(colorSrc.a, colorMsk.a); } + /* DifferenceMask */ case 9u: { result = abs(colorMsk.xyz - colorSrc.xyz * (1.0 - colorMsk.a)); } + default: { result = colorSrc.xyz; } } - let S: vec3f = color.xyz; - let D: vec3f = colorDst.xyz; + let So: f32 = opacity; let Sa: f32 = alpha; let Da: f32 = colorDst.a; - let One: vec3f = vec3(1.0); + let S: vec4f = vec4(result, alpha); + let D: vec4f = colorDst; + let One: vec4f = vec4(1.0); + + var color: vec4f = vec4f(0.0, 0.0, 0.0, 0.0); switch blendMethod { - /* Normal */ case 0u: { color = (Sa * S) + (1.0 - Sa) * D; } + /* 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); } @@ -509,10 +516,10 @@ fn cs_main( @builtin(global_invocation_id) id: vec3u) { /* 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 = (Sa * S) + (1.0 - Sa) * D; } + default: { color = (S * So) + (1.0 - Sa * So) * D; } } - textureStore(imageDst, id.xy, vec4f(color, Sa)); + textureStore(imageDst, id.xy, color); } )";