diff --git a/src/renderer/wg_engine/tvgWgBindGroups.cpp b/src/renderer/wg_engine/tvgWgBindGroups.cpp index 74a30c7f..3fbea8fc 100644 --- a/src/renderer/wg_engine/tvgWgBindGroups.cpp +++ b/src/renderer/wg_engine/tvgWgBindGroups.cpp @@ -35,6 +35,7 @@ WGPUBindGroupLayout WgBindGroupPicture::layout = nullptr; WGPUBindGroupLayout WgBindGroupTexture::layout = nullptr; WGPUBindGroupLayout WgBindGroupTextureStorage::layout = nullptr; WGPUBindGroupLayout WgBindGroupTextureSampled::layout = nullptr; +WGPUBindGroupLayout WgBindGroupTexComposeBlend::layout = nullptr; WGPUBindGroupLayout WgBindGroupOpacity::layout = nullptr; WGPUBindGroupLayout WgBindGroupBlendMethod::layout = nullptr; WGPUBindGroupLayout WgBindGroupCompositeMethod::layout = nullptr; @@ -386,6 +387,46 @@ void WgBindGroupTextureSampled::release() } +WGPUBindGroupLayout WgBindGroupTexComposeBlend::getLayout(WGPUDevice device) +{ + if (layout) return layout; + const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { + makeBindGroupLayoutEntryStorageTexture(0, WGPUStorageTextureAccess_ReadOnly), + makeBindGroupLayoutEntryStorageTexture(1, WGPUStorageTextureAccess_ReadOnly), + makeBindGroupLayoutEntryStorageTexture(2, WGPUStorageTextureAccess_ReadWrite) + }; + layout = createBindGroupLayout(device, bindGroupLayoutEntries, 3); + assert(layout); + return layout; +} + + +void WgBindGroupTexComposeBlend::releaseLayout() +{ + releaseBindGroupLayout(layout); +} + + +void WgBindGroupTexComposeBlend::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexSrc, WGPUTextureView uTexMsk, WGPUTextureView uTexDst) +{ + release(); + const WGPUBindGroupEntry bindGroupEntries[] { + makeBindGroupEntryTextureView(0, uTexSrc), + makeBindGroupEntryTextureView(1, uTexMsk), + makeBindGroupEntryTextureView(2, uTexDst) + }; + mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 3); + assert(mBindGroup); +} + + +void WgBindGroupTexComposeBlend::release() +{ + releaseBindGroup(mBindGroup); +} + + + WGPUBindGroupLayout WgBindGroupOpacity::getLayout(WGPUDevice device) { if (layout) return layout; diff --git a/src/renderer/wg_engine/tvgWgBindGroups.h b/src/renderer/wg_engine/tvgWgBindGroups.h index 921bad35..6f18eb78 100644 --- a/src/renderer/wg_engine/tvgWgBindGroups.h +++ b/src/renderer/wg_engine/tvgWgBindGroups.h @@ -143,6 +143,21 @@ struct WgBindGroupTextureSampled : public WgBindGroup void release(); }; +// @group(0) +struct WgBindGroupTexComposeBlend : public WgBindGroup +{ + static WGPUBindGroupLayout layout; + static WGPUBindGroupLayout getLayout(WGPUDevice device); + static void releaseLayout(); + + void initialize(WGPUDevice device, WGPUQueue queue, + WGPUTextureView uTexSrc, + WGPUTextureView uTexMsk, + WGPUTextureView uTexDst); + void release(); +}; + + // @group(1 or 2) struct WgBindGroupOpacity : public WgBindGroup { diff --git a/src/renderer/wg_engine/tvgWgPipelines.cpp b/src/renderer/wg_engine/tvgWgPipelines.cpp index e14af258..de969815 100644 --- a/src/renderer/wg_engine/tvgWgPipelines.cpp +++ b/src/renderer/wg_engine/tvgWgPipelines.cpp @@ -292,6 +292,28 @@ void WgPipelineCompose::initialize(WGPUDevice device) } +void WgPipelineComposeBlend::initialize(WGPUDevice device) +{ + // bind groups and layouts + WGPUBindGroupLayout bindGroupLayouts[] = { + WgBindGroupTexComposeBlend::getLayout(device), + WgBindGroupCompositeMethod::getLayout(device), + WgBindGroupBlendMethod::getLayout(device), + WgBindGroupOpacity::getLayout(device) + }; + + // sheder source and labels + auto shaderSource = cShaderSource_PipelineComputeComposeBlend; + auto shaderLabel = "The compute shader compose blend"; + auto pipelineLabel = "The compute pipeline compose blend"; + + // allocate all pipeline handles + allocate(device, + bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts), + shaderSource, shaderLabel, pipelineLabel); +} + + void WgPipelineAntiAliasing::initialize(WGPUDevice device) { // bind groups and layouts @@ -330,6 +352,7 @@ void WgPipelines::initialize(WgContext& context) computeClear.initialize(context.device); computeBlend.initialize(context.device); computeCompose.initialize(context.device); + computeComposeBlend.initialize(context.device); computeAntiAliasing.initialize(context.device); // store pipelines to context context.pipelines = this; @@ -338,6 +361,7 @@ void WgPipelines::initialize(WgContext& context) void WgPipelines::release() { + WgBindGroupTexComposeBlend::layout = nullptr; WgBindGroupTextureSampled::releaseLayout(); WgBindGroupTextureStorage::releaseLayout(); WgBindGroupTexture::releaseLayout(); @@ -350,6 +374,7 @@ void WgPipelines::release() WgBindGroupCanvas::releaseLayout(); // compute pipelines computeAntiAliasing.release(); + computeComposeBlend.release(); computeCompose.release(); computeBlend.release(); computeClear.release(); diff --git a/src/renderer/wg_engine/tvgWgPipelines.h b/src/renderer/wg_engine/tvgWgPipelines.h index 61e58fba..8d477dd1 100644 --- a/src/renderer/wg_engine/tvgWgPipelines.h +++ b/src/renderer/wg_engine/tvgWgPipelines.h @@ -145,6 +145,20 @@ struct WgPipelineCompose: public WgComputePipeline }; +struct WgPipelineComposeBlend: public WgComputePipeline +{ + void initialize(WGPUDevice device) override; + void use(WGPUComputePassEncoder encoder, WgBindGroupTexComposeBlend& groupTexs, WgBindGroupCompositeMethod& groupComposeMethod, WgBindGroupBlendMethod& groupBlendMethod, WgBindGroupOpacity& groupOpacity) + { + set(encoder); + groupTexs.set(encoder, 0); + groupComposeMethod.set(encoder, 1); + groupBlendMethod.set(encoder, 2); + groupOpacity.set(encoder, 3); + } +}; + + struct WgPipelineAntiAliasing: public WgComputePipeline { void initialize(WGPUDevice device) override; @@ -174,6 +188,7 @@ struct WgPipelines WgPipelineClear computeClear; WgPipelineBlend computeBlend; WgPipelineCompose computeCompose; + WgPipelineComposeBlend computeComposeBlend; WgPipelineAntiAliasing computeAntiAliasing; void initialize(WgContext& context); diff --git a/src/renderer/wg_engine/tvgWgRenderTarget.cpp b/src/renderer/wg_engine/tvgWgRenderTarget.cpp index 9ac5e88c..f516bd66 100644 --- a/src/renderer/wg_engine/tvgWgRenderTarget.cpp +++ b/src/renderer/wg_engine/tvgWgRenderTarget.cpp @@ -185,6 +185,28 @@ void WgRenderStorage::compose(WGPUCommandEncoder commandEncoder, WgRenderStorage }; +void WgRenderStorage::composeBlend( + WgContext& context, + WGPUCommandEncoder commandEncoder, + WgRenderStorage* texSrc, + WgRenderStorage* texMsk, + WgBindGroupCompositeMethod* composeMethod, + WgBindGroupBlendMethod* blendMethod, + WgBindGroupOpacity* opacity) +{ + assert(commandEncoder); + assert(texSrc); + assert(texMsk); + WgBindGroupTexComposeBlend composeBlend; + composeBlend.initialize(context.device, context.queue, texSrc->texViewColor, texMsk->texViewColor, texViewColor); + WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder); + mPipelines->computeComposeBlend.use(computePassEncoder, composeBlend, *composeMethod, *blendMethod, *opacity); + dispatchWorkgroups(computePassEncoder); + endComputePass(computePassEncoder); + composeBlend.release(); +} + + void WgRenderStorage::antialias(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc) { assert(commandEncoder); diff --git a/src/renderer/wg_engine/tvgWgRenderTarget.h b/src/renderer/wg_engine/tvgWgRenderTarget.h index 219a5dbe..2670eaa4 100644 --- a/src/renderer/wg_engine/tvgWgRenderTarget.h +++ b/src/renderer/wg_engine/tvgWgRenderTarget.h @@ -54,6 +54,14 @@ public: void clear(WGPUCommandEncoder commandEncoder); void blend(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc, WgBindGroupBlendMethod* blendMethod); void compose(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetMsk, WgBindGroupCompositeMethod* composeMethod, WgBindGroupOpacity* opacity); + void composeBlend( + WgContext& context, + WGPUCommandEncoder commandEncoder, + WgRenderStorage* texMsk, + WgRenderStorage* texSrc, + WgBindGroupCompositeMethod* composeMethod, + WgBindGroupBlendMethod* blendMethod, + WgBindGroupOpacity* opacity); void antialias(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc); private: void drawShape(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType); diff --git a/src/renderer/wg_engine/tvgWgRenderer.cpp b/src/renderer/wg_engine/tvgWgRenderer.cpp index 79680a56..408aa7a8 100644 --- a/src/renderer/wg_engine/tvgWgRenderer.cpp +++ b/src/renderer/wg_engine/tvgWgRenderer.cpp @@ -392,20 +392,24 @@ bool WgRenderer::endComposite(TVG_UNUSED Compositor* cmp) } else { // end current render pass mRenderStorageStack.last()->endRenderPass(); - // get two last render targets + + // get source, mask and destination render storages WgRenderStorage* renderStorageSrc = mRenderStorageStack.last(); mRenderStorageStack.pop(); WgRenderStorage* renderStorageMsk = mRenderStorageStack.last(); mRenderStorageStack.pop(); - - // compose shape and mask - WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, cmp->opacity); - WgBindGroupCompositeMethod* composeMethod = mCompositeMethodPool.allocate(mContext, cmp->method); - renderStorageSrc->compose(mCommandEncoder, renderStorageMsk, composeMethod, opacity); + WgRenderStorage* renderStorageDst = mRenderStorageStack.last(); - // blent scene to current render storage + // get compose, blend and opacity settings + WgBindGroupCompositeMethod* composeMethod = mCompositeMethodPool.allocate(mContext, cmp->method); WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod); - mRenderStorageStack.last()->blend(mCommandEncoder, renderStorageSrc, blendMethod); + WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, cmp->opacity); + + // compose and blend + // dest = blend(dest, compose(src, msk, composeMethod), blendMethod, opacity) + renderStorageDst->composeBlend(mContext, mCommandEncoder, + renderStorageSrc, renderStorageMsk, + composeMethod, 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 50bd1deb..00841183 100644 --- a/src/renderer/wg_engine/tvgWgShaderSrc.cpp +++ b/src/renderer/wg_engine/tvgWgShaderSrc.cpp @@ -453,6 +453,68 @@ fn cs_main( @builtin(global_invocation_id) id: vec3u) { } )"; +// pipeline shader modules compose blend +const char* cShaderSource_PipelineComputeComposeBlend = R"( +@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; +@group(1) @binding(0) var composeMethod : u32; +@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 colorSrc = textureLoad(imageSrc, id.xy); + let colorMsk = textureLoad(imageMsk, id.xy); + let colorDst = textureLoad(imageDst, id.xy); + + var color: vec3f = colorSrc.xyz; + var alpha: f32 = colorMsk.a; + let luma: vec3f = vec3f(0.299, 0.587, 0.114); + switch composeMethod { + /* None */ case 0u: { color = colorSrc.xyz; } + /* ClipPath */ case 1u: { if (colorMsk.a == 0) { alpha = 0.0; }; } + /* AlphaMask */ case 2u: { color = mix(colorMsk.xyz, colorSrc.xyz, colorSrc.a * colorMsk.b); } + /* InvAlphaMask */ case 3u: { color = mix(colorSrc.xyz, colorMsk.xyz, colorSrc.a * colorMsk.b); alpha = 1.0 - colorMsk.b; } + /* LumaMask */ case 4u: { color = colorSrc.xyz * dot(colorMsk.xyz, luma); } + /* InvLumaMask */ case 5u: { color = colorSrc.xyz * (1.0 - dot(colorMsk.xyz, luma)); alpha = 1.0 - colorMsk.b; } + /* 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; } + } + + let S: vec3f = color.xyz; + let D: vec3f = colorDst.xyz; + let Sa: f32 = alpha * opacity; + let Da: f32 = colorDst.a; + let One: vec3f = vec3(1.0); + switch blendMethod { + /* Normal */ case 0u: { color = (Sa * S) + (1.0 - Sa) * 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 = S * 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 = (Sa * S) + (1.0 - Sa) * D; } + } + + textureStore(imageDst, id.xy, vec4f(color, Sa)); +} +)"; + // pipeline shader modules anti-aliasing const char* cShaderSource_PipelineComputeAntiAlias = R"( @group(0) @binding(0) var imageSrc : texture_storage_2d; diff --git a/src/renderer/wg_engine/tvgWgShaderSrc.h b/src/renderer/wg_engine/tvgWgShaderSrc.h index 7eec4b62..60d7c83b 100644 --- a/src/renderer/wg_engine/tvgWgShaderSrc.h +++ b/src/renderer/wg_engine/tvgWgShaderSrc.h @@ -44,6 +44,7 @@ extern const char* cShaderSource_PipelineImage; extern const char* cShaderSource_PipelineComputeClear; extern const char* cShaderSource_PipelineComputeBlend; extern const char* cShaderSource_PipelineComputeCompose; +extern const char* cShaderSource_PipelineComputeComposeBlend; extern const char* cShaderSource_PipelineComputeAntiAlias; #endif // _TVG_WG_SHADER_SRC_H_