diff --git a/src/renderer/wg_engine/tvgWgBindGroups.cpp b/src/renderer/wg_engine/tvgWgBindGroups.cpp index 891bead4..07e76697 100644 --- a/src/renderer/wg_engine/tvgWgBindGroups.cpp +++ b/src/renderer/wg_engine/tvgWgBindGroups.cpp @@ -36,11 +36,15 @@ WGPUBindGroupLayout WgBindGroupTexture::layout = nullptr; WGPUBindGroupLayout WgBindGroupTextureStorageRgba::layout = nullptr; WGPUBindGroupLayout WgBindGroupTextureStorageBgra::layout = nullptr; WGPUBindGroupLayout WgBindGroupTextureSampled::layout = nullptr; +WGPUBindGroupLayout WgBindGroupTexMaskCompose::layout = nullptr; WGPUBindGroupLayout WgBindGroupTexComposeBlend::layout = nullptr; WGPUBindGroupLayout WgBindGroupOpacity::layout = nullptr; WGPUBindGroupLayout WgBindGroupBlendMethod::layout = nullptr; WGPUBindGroupLayout WgBindGroupCompositeMethod::layout = nullptr; +/////////////////////////////////////////////////////////////////////////////// +// WgBindGroupCanvas +/////////////////////////////////////////////////////////////////////////////// WGPUBindGroupLayout WgBindGroupCanvas::getLayout(WGPUDevice device) { @@ -78,6 +82,9 @@ void WgBindGroupCanvas::release() releaseBuffer(uBufferViewMat); } +/////////////////////////////////////////////////////////////////////////////// +// WgBindGroupPaint +/////////////////////////////////////////////////////////////////////////////// WGPUBindGroupLayout WgBindGroupPaint::getLayout(WGPUDevice device) { @@ -123,6 +130,9 @@ void WgBindGroupPaint::release() releaseBuffer(uBufferModelMat); } +/////////////////////////////////////////////////////////////////////////////// +// WgBindGroupSolidColor +/////////////////////////////////////////////////////////////////////////////// WGPUBindGroupLayout WgBindGroupSolidColor::getLayout(WGPUDevice device) { @@ -163,6 +173,9 @@ void WgBindGroupSolidColor::release() releaseBuffer(uBufferSolidColor); } +/////////////////////////////////////////////////////////////////////////////// +// WgBindGroupLinearGradient +/////////////////////////////////////////////////////////////////////////////// WGPUBindGroupLayout WgBindGroupLinearGradient::getLayout(WGPUDevice device) { @@ -203,6 +216,9 @@ void WgBindGroupLinearGradient::release() releaseBuffer(uBufferLinearGradient); } +/////////////////////////////////////////////////////////////////////////////// +// WgBindGroupRadialGradient +/////////////////////////////////////////////////////////////////////////////// WGPUBindGroupLayout WgBindGroupRadialGradient::getLayout(WGPUDevice device) { @@ -243,6 +259,9 @@ void WgBindGroupRadialGradient::release() releaseBindGroup(mBindGroup); } +/////////////////////////////////////////////////////////////////////////////// +// WgBindGroupPicture +/////////////////////////////////////////////////////////////////////////////// WGPUBindGroupLayout WgBindGroupPicture::getLayout(WGPUDevice device) { @@ -280,6 +299,9 @@ void WgBindGroupPicture::release() releaseBindGroup(mBindGroup); } +/////////////////////////////////////////////////////////////////////////////// +// WgBindGroupTexture +/////////////////////////////////////////////////////////////////////////////// WGPUBindGroupLayout WgBindGroupTexture::getLayout(WGPUDevice device) { @@ -315,6 +337,9 @@ void WgBindGroupTexture::release() releaseBindGroup(mBindGroup); } +/////////////////////////////////////////////////////////////////////////////// +// WgBindGroupTextureStorageRgba +/////////////////////////////////////////////////////////////////////////////// WGPUBindGroupLayout WgBindGroupTextureStorageRgba::getLayout(WGPUDevice device) { @@ -350,6 +375,9 @@ void WgBindGroupTextureStorageRgba::release() releaseBindGroup(mBindGroup); } +/////////////////////////////////////////////////////////////////////////////// +// WgBindGroupTextureStorageBgra +/////////////////////////////////////////////////////////////////////////////// WGPUBindGroupLayout WgBindGroupTextureStorageBgra::getLayout(WGPUDevice device) { @@ -385,6 +413,9 @@ void WgBindGroupTextureStorageBgra::release() releaseBindGroup(mBindGroup); } +/////////////////////////////////////////////////////////////////////////////// +// WgBindGroupTextureSampled +/////////////////////////////////////////////////////////////////////////////// WGPUBindGroupLayout WgBindGroupTextureSampled::getLayout(WGPUDevice device) { @@ -422,6 +453,9 @@ void WgBindGroupTextureSampled::release() releaseBindGroup(mBindGroup); } +/////////////////////////////////////////////////////////////////////////////// +// WgBindGroupTexComposeBlend +/////////////////////////////////////////////////////////////////////////////// WGPUBindGroupLayout WgBindGroupTexComposeBlend::getLayout(WGPUDevice device) { @@ -461,7 +495,49 @@ void WgBindGroupTexComposeBlend::release() releaseBindGroup(mBindGroup); } +/////////////////////////////////////////////////////////////////////////////// +// WgBindGroupTexMaskCompose +/////////////////////////////////////////////////////////////////////////////// +WGPUBindGroupLayout WgBindGroupTexMaskCompose::getLayout(WGPUDevice device) +{ + if (layout) return layout; + const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { + makeBindGroupLayoutEntryStorage(0, WGPUStorageTextureAccess_ReadOnly, WGPUTextureFormat_RGBA8Unorm), + makeBindGroupLayoutEntryStorage(1, WGPUStorageTextureAccess_ReadWrite, WGPUTextureFormat_RGBA8Unorm) + }; + layout = createBindGroupLayout(device, bindGroupLayoutEntries, 2); + assert(layout); + return layout; +} + + +void WgBindGroupTexMaskCompose::releaseLayout() +{ + releaseBindGroupLayout(layout); +} + + +void WgBindGroupTexMaskCompose::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexMsk0, WGPUTextureView uTexMsk1) +{ + release(); + const WGPUBindGroupEntry bindGroupEntries[] { + makeBindGroupEntryTextureView(0, uTexMsk0), + makeBindGroupEntryTextureView(1, uTexMsk1), + }; + mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 2); + assert(mBindGroup); +} + + +void WgBindGroupTexMaskCompose::release() +{ + releaseBindGroup(mBindGroup); +} + +/////////////////////////////////////////////////////////////////////////////// +// WgBindGroupOpacity +/////////////////////////////////////////////////////////////////////////////// WGPUBindGroupLayout WgBindGroupOpacity::getLayout(WGPUDevice device) { @@ -500,6 +576,9 @@ void WgBindGroupOpacity::release() releaseBindGroup(mBindGroup); } +/////////////////////////////////////////////////////////////////////////////// +// WgBindGroupBlendMethod +/////////////////////////////////////////////////////////////////////////////// WGPUBindGroupLayout WgBindGroupBlendMethod::getLayout(WGPUDevice device) { @@ -538,6 +617,9 @@ void WgBindGroupBlendMethod::release() releaseBindGroup(mBindGroup); } +/////////////////////////////////////////////////////////////////////////////// +// WgBindGroupCompositeMethod +/////////////////////////////////////////////////////////////////////////////// WGPUBindGroupLayout WgBindGroupCompositeMethod::getLayout(WGPUDevice device) { diff --git a/src/renderer/wg_engine/tvgWgBindGroups.h b/src/renderer/wg_engine/tvgWgBindGroups.h index 2e01651b..0d9b33e1 100644 --- a/src/renderer/wg_engine/tvgWgBindGroups.h +++ b/src/renderer/wg_engine/tvgWgBindGroups.h @@ -170,6 +170,20 @@ struct WgBindGroupTexComposeBlend : public WgBindGroup }; +// @group(0) +struct WgBindGroupTexMaskCompose : public WgBindGroup +{ + static WGPUBindGroupLayout layout; + static WGPUBindGroupLayout getLayout(WGPUDevice device); + static void releaseLayout(); + + void initialize(WGPUDevice device, WGPUQueue queue, + WGPUTextureView uTexMsk0, + WGPUTextureView uTexMsk1); + 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 567dd0bf..4e0ce4de 100644 --- a/src/renderer/wg_engine/tvgWgPipelines.cpp +++ b/src/renderer/wg_engine/tvgWgPipelines.cpp @@ -49,7 +49,7 @@ void WgPipelineFillShapeWinding::initialize(WGPUDevice device) WGPUCompareFunction stencilFunctionBack = WGPUCompareFunction_Always; WGPUStencilOperation stencilOperationBack = WGPUStencilOperation_DecrementWrap; - // sheder source and labels + // shader source and labels auto shaderSource = cShaderSource_PipelineFill; auto shaderLabel = "The shader fill"; auto pipelineLabel = "The render pipeline fill shape winding"; @@ -83,7 +83,7 @@ void WgPipelineFillShapeEvenOdd::initialize(WGPUDevice device) WGPUCompareFunction stencilFunctionBack = WGPUCompareFunction_Always; WGPUStencilOperation stencilOperationBack = WGPUStencilOperation_Invert; - // sheder source and labels + // shader source and labels auto shaderSource = cShaderSource_PipelineFill; auto shaderLabel = "The shader fill"; auto pipelineLabel = "The render pipeline fill shape Even Odd"; @@ -115,7 +115,7 @@ void WgPipelineFillStroke::initialize(WGPUDevice device) WGPUCompareFunction stencilFunction = WGPUCompareFunction_Always; WGPUStencilOperation stencilOperation = WGPUStencilOperation_Replace; - // sheder source and labels + // shader source and labels auto shaderSource = cShaderSource_PipelineFill; auto shaderLabel = "The shader fill"; auto pipelineLabel = "The render pipeline fill stroke"; @@ -129,6 +129,38 @@ void WgPipelineFillStroke::initialize(WGPUDevice device) } +void WgPipelineClipMask::initialize(WGPUDevice device) +{ + // vertex and buffers settings + WGPUVertexAttribute vertexAttributesPos = { WGPUVertexFormat_Float32x2, sizeof(float) * 0, 0 }; + WGPUVertexBufferLayout vertexBufferLayouts[] = { + makeVertexBufferLayout(&vertexAttributesPos, 1, sizeof(float) * 2) + }; + + // bind groups + WGPUBindGroupLayout bindGroupLayouts[] = { + WgBindGroupCanvas::getLayout(device), + WgBindGroupPaint::getLayout(device) + }; + + // stencil function + WGPUCompareFunction stencilFunction = WGPUCompareFunction_NotEqual; + WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero; + + // shader source and labels + auto shaderSource = cShaderSource_PipelineFill; + auto shaderLabel = "The shader fill"; + auto pipelineLabel = "The render pipeline clip mask"; + + // allocate all pipeline handles + allocate(device, WgPipelineBlendType::SrcOver, WGPUColorWriteMask_All, + vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts), + bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts), + stencilFunction, stencilOperation, stencilFunction, stencilOperation, + shaderSource, shaderLabel, pipelineLabel); +} + + void WgPipelineSolid::initialize(WGPUDevice device, WgPipelineBlendType blendType) { // vertex and buffers settings @@ -148,7 +180,7 @@ void WgPipelineSolid::initialize(WGPUDevice device, WgPipelineBlendType blendTyp WGPUCompareFunction stencilFunction = WGPUCompareFunction_NotEqual; WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero; - // sheder source and labels + // shader source and labels auto shaderSource = cShaderSource_PipelineSolid; auto shaderLabel = "The shader solid color"; auto pipelineLabel = "The render pipeline solid color"; @@ -181,7 +213,7 @@ void WgPipelineLinear::initialize(WGPUDevice device, WgPipelineBlendType blendTy WGPUCompareFunction stencilFunction = WGPUCompareFunction_NotEqual; WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero; - // sheder source and labels + // shader source and labels auto shaderSource = cShaderSource_PipelineLinear; auto shaderLabel = "The shader linear gradient"; auto pipelineLabel = "The render pipeline linear gradient"; @@ -214,7 +246,7 @@ void WgPipelineRadial::initialize(WGPUDevice device, WgPipelineBlendType blendTy WGPUCompareFunction stencilFunction = WGPUCompareFunction_NotEqual; WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero; - // sheder source and labels + // shader source and labels auto shaderSource = cShaderSource_PipelineRadial; auto shaderLabel = "The shader radial gradient"; auto pipelineLabel = "The render pipeline radial gradient"; @@ -249,7 +281,7 @@ void WgPipelineImage::initialize(WGPUDevice device, WgPipelineBlendType blendTyp WGPUCompareFunction stencilFunction = WGPUCompareFunction_Always; WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero; - // sheder source and labels + // shader source and labels auto shaderSource = cShaderSource_PipelineImage; auto shaderLabel = "The shader image"; auto pipelineLabel = "The render pipeline image"; @@ -273,7 +305,7 @@ void WgPipelineClear::initialize(WGPUDevice device) WgBindGroupTextureStorageRgba::getLayout(device) }; - // sheder source and labels + // shader source and labels auto shaderSource = cShaderSource_PipelineComputeClear; auto shaderLabel = "The compute shader clear"; auto pipelineLabel = "The compute pipeline clear"; @@ -295,7 +327,7 @@ void WgPipelineBlend::initialize(WGPUDevice device, const char *shaderSource) WgBindGroupOpacity::getLayout(device) }; - // sheder source and labels + // shader source and labels auto shaderLabel = "The compute shader blend"; auto pipelineLabel = "The compute pipeline blend"; @@ -306,7 +338,46 @@ void WgPipelineBlend::initialize(WGPUDevice device, const char *shaderSource) } -void WgPipelineComposeBlend::initialize(WGPUDevice device) +void WgPipelineBlendMask::initialize(WGPUDevice device, const char *shaderSource) +{ + // bind groups and layouts + WGPUBindGroupLayout bindGroupLayouts[] = { + WgBindGroupTexComposeBlend::getLayout(device), + WgBindGroupBlendMethod::getLayout(device), + WgBindGroupOpacity::getLayout(device) + }; + + // shader source and labels + auto shaderLabel = "The compute shader blend mask"; + auto pipelineLabel = "The compute pipeline blend mask"; + + // allocate all pipeline handles + allocate(device, + bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts), + shaderSource, shaderLabel, pipelineLabel); +} + + +void WgPipelineMaskCompose::initialize(WGPUDevice device) +{ + // bind groups and layouts + WGPUBindGroupLayout bindGroupLayouts[] = { + WgBindGroupTexMaskCompose::getLayout(device), + }; + + // shader source and labels + auto shaderSource = cShaderSource_PipelineComputeMaskCompose; + auto shaderLabel = "The compute shader mask compose"; + auto pipelineLabel = "The compute pipeline mask compose"; + + // allocate all pipeline handles + allocate(device, + bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts), + shaderSource, shaderLabel, pipelineLabel); +} + + +void WgPipelineCompose::initialize(WGPUDevice device) { // bind groups and layouts WGPUBindGroupLayout bindGroupLayouts[] = { @@ -316,8 +387,8 @@ void WgPipelineComposeBlend::initialize(WGPUDevice device) WgBindGroupOpacity::getLayout(device) }; - // sheder source and labels - auto shaderSource = cShaderSource_PipelineComputeComposeBlend; + // shader source and labels + auto shaderSource = cShaderSource_PipelineComputeCompose; auto shaderLabel = "The compute shader compose blend"; auto pipelineLabel = "The compute pipeline compose blend"; @@ -336,7 +407,7 @@ void WgPipelineAntiAliasing::initialize(WGPUDevice device) WgBindGroupTextureStorageBgra::getLayout(device) }; - // sheder source and labels + // shader source and labels auto shaderSource = cShaderSource_PipelineComputeAntiAlias; auto shaderLabel = "The compute shader anti-aliasing"; auto pipelineLabel = "The compute pipeline anti-aliasing"; @@ -357,6 +428,7 @@ void WgPipelines::initialize(WgContext& context) fillShapeWinding.initialize(context.device); fillShapeEvenOdd.initialize(context.device); fillStroke.initialize(context.device); + clipMask.initialize(context.device); 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); @@ -368,7 +440,11 @@ void WgPipelines::initialize(WgContext& context) computeBlendSolid.initialize(context.device, cShaderSource_PipelineComputeBlendSolid); computeBlendGradient.initialize(context.device, cShaderSource_PipelineComputeBlendGradient); computeBlendImage.initialize(context.device, cShaderSource_PipelineComputeBlendImage); - computeComposeBlend.initialize(context.device); + computeBlendSolidMask.initialize(context.device, cShaderSource_PipelineComputeBlendSolidMask); + computeBlendGradientMask.initialize(context.device, cShaderSource_PipelineComputeBlendGradientMask); + computeBlendImageMask.initialize(context.device, cShaderSource_PipelineComputeBlendImageMask); + computeMaskCompose.initialize(context.device); + computeCompose.initialize(context.device); computeAntiAliasing.initialize(context.device); // store pipelines to context context.pipelines = this; @@ -377,7 +453,8 @@ void WgPipelines::initialize(WgContext& context) void WgPipelines::release() { - WgBindGroupTexComposeBlend::layout = nullptr; + WgBindGroupTexMaskCompose::releaseLayout(); + WgBindGroupTexComposeBlend::releaseLayout(); WgBindGroupTextureSampled::releaseLayout(); WgBindGroupTextureStorageBgra::releaseLayout(); WgBindGroupTextureStorageRgba::releaseLayout(); @@ -391,7 +468,11 @@ void WgPipelines::release() WgBindGroupCanvas::releaseLayout(); // compute pipelines computeAntiAliasing.release(); - computeComposeBlend.release(); + computeCompose.release(); + computeMaskCompose.release(); + computeBlendImageMask.release(); + computeBlendGradientMask.release(); + computeBlendSolidMask.release(); computeBlendImage.release(); computeBlendGradient.release(); computeBlendSolid.release(); @@ -403,6 +484,7 @@ void WgPipelines::release() linear[type].release(); solid[type].release(); } + clipMask.release(); fillStroke.release(); fillShapeEvenOdd.release(); fillShapeWinding.release(); diff --git a/src/renderer/wg_engine/tvgWgPipelines.h b/src/renderer/wg_engine/tvgWgPipelines.h index c88a7892..2a94a9cd 100644 --- a/src/renderer/wg_engine/tvgWgPipelines.h +++ b/src/renderer/wg_engine/tvgWgPipelines.h @@ -62,6 +62,17 @@ struct WgPipelineFillStroke: public WgRenderPipeline } }; +struct WgPipelineClipMask: public WgRenderPipeline +{ + void initialize(WGPUDevice device) override; + void use(WGPURenderPassEncoder encoder, WgBindGroupCanvas& groupCanvas,WgBindGroupPaint& groupPaint) + { + set(encoder); + groupCanvas.set(encoder, 0); + groupPaint.set(encoder, 1); + } +}; + struct WgPipelineSolid: public WgRenderPipeline { void initialize(WGPUDevice device) override {} @@ -144,7 +155,32 @@ struct WgPipelineBlend: public WgComputePipeline }; -struct WgPipelineComposeBlend: public WgComputePipeline +struct WgPipelineBlendMask: public WgComputePipeline +{ + void initialize(WGPUDevice device) override { assert(false); }; + void initialize(WGPUDevice device, const char *shaderSource); + void use(WGPUComputePassEncoder encoder, WgBindGroupTexComposeBlend& groupTexs, WgBindGroupBlendMethod& blendMethod, WgBindGroupOpacity& groupOpacity) + { + set(encoder); + groupTexs.set(encoder, 0); + blendMethod.set(encoder, 1); + groupOpacity.set(encoder, 2); + } +}; + + +struct WgPipelineMaskCompose: public WgComputePipeline +{ + void initialize(WGPUDevice device) override; + void use(WGPUComputePassEncoder encoder, WgBindGroupTexMaskCompose& groupTexs) + { + set(encoder); + groupTexs.set(encoder, 0); + } +}; + + +struct WgPipelineCompose: public WgComputePipeline { void initialize(WGPUDevice device) override; void use(WGPUComputePassEncoder encoder, WgBindGroupTexComposeBlend& groupTexs, WgBindGroupCompositeMethod& groupComposeMethod, WgBindGroupBlendMethod& groupBlendMethod, WgBindGroupOpacity& groupOpacity) @@ -180,6 +216,7 @@ struct WgPipelines WgPipelineFillShapeEvenOdd fillShapeEvenOdd; WgPipelineFillStroke fillStroke; // fill pipelines + WgPipelineClipMask clipMask; WgPipelineSolid solid[3]; WgPipelineLinear linear[3]; WgPipelineRadial radial[3]; @@ -189,7 +226,11 @@ struct WgPipelines WgPipelineBlend computeBlendSolid; WgPipelineBlend computeBlendGradient; WgPipelineBlend computeBlendImage; - WgPipelineComposeBlend computeComposeBlend; + WgPipelineBlendMask computeBlendSolidMask; + WgPipelineBlendMask computeBlendGradientMask; + WgPipelineBlendMask computeBlendImageMask; + WgPipelineMaskCompose computeMaskCompose; + WgPipelineCompose computeCompose; WgPipelineAntiAliasing computeAntiAliasing; void initialize(WgContext& context); diff --git a/src/renderer/wg_engine/tvgWgRenderData.cpp b/src/renderer/wg_engine/tvgWgRenderData.cpp index 0b9e94b9..4a995619 100644 --- a/src/renderer/wg_engine/tvgWgRenderData.cpp +++ b/src/renderer/wg_engine/tvgWgRenderData.cpp @@ -262,8 +262,17 @@ void WgRenderSettings::release(WgContext& context) void WgRenderDataPaint::release(WgContext& context) { bindGroupPaint.release(); + clips.clear(); }; + +void WgRenderDataPaint::updateClips(tvg::Array &clips) { + this->clips.clear(); + for (uint32_t i = 0; i < clips.count; i++) + if (clips[i]) + this->clips.push((WgRenderDataPaint*)clips[i]); +} + //*********************************************************************** // WgRenderDataShape //*********************************************************************** @@ -371,6 +380,7 @@ void WgRenderDataShape::releaseMeshes(WgContext &context) meshGroupShapes.release(context); pMin = {FLT_MAX, FLT_MAX}; pMax = {0.0f, 0.0f}; + clips.clear(); } @@ -407,6 +417,7 @@ void WgRenderDataShapePool::free(WgContext& context, WgRenderDataShape* dataShap dataShape->meshGroupShapesBBox.release(context); dataShape->meshGroupStrokes.release(context); dataShape->meshGroupStrokesBBox.release(context); + dataShape->clips.clear(); mPool.push(dataShape); } diff --git a/src/renderer/wg_engine/tvgWgRenderData.h b/src/renderer/wg_engine/tvgWgRenderData.h index 84ea8aaf..f41d5642 100644 --- a/src/renderer/wg_engine/tvgWgRenderData.h +++ b/src/renderer/wg_engine/tvgWgRenderData.h @@ -88,10 +88,13 @@ struct WgRenderDataPaint WgBindGroupPaint bindGroupPaint{}; RenderRegion viewport{}; float opacity{}; + Array clips; virtual ~WgRenderDataPaint() {}; virtual void release(WgContext& context); virtual Type type() { return Type::Undefined; }; + + void updateClips(tvg::Array &clips); }; struct WgRenderDataShape: public WgRenderDataPaint diff --git a/src/renderer/wg_engine/tvgWgRenderTarget.cpp b/src/renderer/wg_engine/tvgWgRenderTarget.cpp index a936be47..21314580 100644 --- a/src/renderer/wg_engine/tvgWgRenderTarget.cpp +++ b/src/renderer/wg_engine/tvgWgRenderTarget.cpp @@ -85,7 +85,6 @@ void WgRenderStorage::renderShape(WgContext& context, WgRenderDataShape* renderD { assert(renderData); assert(mRenderPassEncoder); - // draw strokes if (renderData->strokeFirst) { drawStroke(context, renderData, blendType); drawShape(context, renderData, blendType); @@ -102,12 +101,21 @@ void WgRenderStorage::renderPicture(WgContext& context, WgRenderDataPicture* ren assert(mRenderPassEncoder); uint8_t blend = (uint8_t)blendType; auto& vp = renderData->viewport; + if ((vp.w <= 0) || (vp.h <= 0)) return; wgpuRenderPassEncoderSetScissorRect(mRenderPassEncoder, vp.x * samples, vp.y * samples, vp.w * samples, vp.h * samples); wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 0); mPipelines->image[blend].use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint, renderData->bindGroupPicture); renderData->meshData.drawImage(context, mRenderPassEncoder); } +void WgRenderStorage::renderClipPath(WgContext& context, WgRenderDataPaint* renderData) { + assert(renderData); + assert(mRenderPassEncoder); + if (renderData->type() == Type::Shape) + drawShapeClipPath(context, (WgRenderDataShape*)renderData); + else if (renderData->type() == Type::Picture) + drawPictureClipPath(context, (WgRenderDataPicture*)renderData); +} void WgRenderStorage::drawShape(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType) { @@ -115,10 +123,12 @@ void WgRenderStorage::drawShape(WgContext& context, WgRenderDataShape* renderDat assert(mRenderPassEncoder); assert(renderData->meshGroupShapes.meshes.count == renderData->meshGroupShapesBBox.meshes.count); if (renderData->renderSettingsShape.skip) return; + // apply viewport auto& vp = renderData->viewport; + if ((vp.w <= 0) || (vp.h <= 0)) return; wgpuRenderPassEncoderSetScissorRect(mRenderPassEncoder, vp.x * samples, vp.y * samples, vp.w * samples, vp.h * samples); - wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 0); // setup fill rule + wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 0); if (renderData->fillRule == FillRule::Winding) mPipelines->fillShapeWinding.use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint); else @@ -145,7 +155,9 @@ void WgRenderStorage::drawStroke(WgContext& context, WgRenderDataShape* renderDa assert(mRenderPassEncoder); assert(renderData->meshGroupStrokes.meshes.count == renderData->meshGroupStrokesBBox.meshes.count); if (renderData->renderSettingsStroke.skip) return; + // apply viewport auto& vp = renderData->viewport; + if ((vp.w <= 0) || (vp.h <= 0)) return; wgpuRenderPassEncoderSetScissorRect(mRenderPassEncoder, vp.x * samples, vp.y * samples, vp.w * samples, vp.h * samples); // draw stroke geometry uint8_t blend = (uint8_t)blendType; @@ -169,6 +181,35 @@ void WgRenderStorage::drawStroke(WgContext& context, WgRenderDataShape* renderDa } +void WgRenderStorage::drawShapeClipPath(WgContext& context, WgRenderDataShape* renderData) { + assert(renderData); + assert(renderData->type() == Type::Shape); + assert(renderData->meshGroupShapes.meshes.count == renderData->meshGroupShapesBBox.meshes.count); + // draw shape geometry + wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 0); + // setup fill rule + if (renderData->fillRule == FillRule::Winding) + mPipelines->fillShapeWinding.use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint); + else + mPipelines->fillShapeEvenOdd.use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint); + // draw to stencil (first pass) + for (uint32_t i = 0; i < renderData->meshGroupShapes.meshes.count; i++) + renderData->meshGroupShapes.meshes[i]->drawFan(context, mRenderPassEncoder); + // fill shape geometry (second pass) + mPipelines->clipMask.use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint); + renderData->meshDataBBox.drawFan(context, mRenderPassEncoder); +} + +void WgRenderStorage::drawPictureClipPath(WgContext& context, WgRenderDataPicture* renderData) { + assert(renderData); + assert(renderData->type() == Type::Picture); + assert(mRenderPassEncoder); + wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 0); + mPipelines->clipMask.use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint); + renderData->meshData.drawImage(context, mRenderPassEncoder); +} + + void WgRenderStorage::clear(WGPUCommandEncoder commandEncoder) { assert(commandEncoder); @@ -179,7 +220,14 @@ void WgRenderStorage::clear(WGPUCommandEncoder commandEncoder) } -void WgRenderStorage::blend(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc, WgPipelineBlend* pipeline, WgBindGroupBlendMethod* blendMethod, WgBindGroupOpacity* opacity) { +void WgRenderStorage::blend( + WgContext& context, + WGPUCommandEncoder commandEncoder, + WgPipelineBlend* pipeline, + WgRenderStorage* targetSrc, + WgBindGroupBlendMethod* blendMethod, + WgBindGroupOpacity* opacity) +{ assert(commandEncoder); assert(targetSrc); assert(pipeline); @@ -190,7 +238,46 @@ void WgRenderStorage::blend(WGPUCommandEncoder commandEncoder, WgRenderStorage* } -void WgRenderStorage::composeBlend( +void WgRenderStorage::blendMask( + WgContext& context, + WGPUCommandEncoder commandEncoder, + WgPipelineBlendMask* pipeline, + WgRenderStorage* texMsk, + WgRenderStorage* texSrc, + 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); + pipeline->use(computePassEncoder, composeBlend, *blendMethod, *opacity); + dispatchWorkgroups(computePassEncoder); + endComputePass(computePassEncoder); + composeBlend.release(); +}; + + +void WgRenderStorage::maskCompose( + WgContext& context, + WGPUCommandEncoder commandEncoder, + WgRenderStorage* texMsk0) +{ + assert(commandEncoder); + assert(texMsk0); + WgBindGroupTexMaskCompose maskCompose; + maskCompose.initialize(context.device, context.queue, texMsk0->texViewColor, texViewColor); + WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder); + mPipelines->computeMaskCompose.use(computePassEncoder, maskCompose); + dispatchWorkgroups(computePassEncoder); + endComputePass(computePassEncoder); + maskCompose.release(); +} + + +void WgRenderStorage::compose( WgContext& context, WGPUCommandEncoder commandEncoder, WgRenderStorage* texSrc, @@ -205,7 +292,7 @@ void WgRenderStorage::composeBlend( 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); + mPipelines->computeCompose.use(computePassEncoder, composeBlend, *composeMethod, *blendMethod, *opacity); dispatchWorkgroups(computePassEncoder); endComputePass(computePassEncoder); composeBlend.release(); diff --git a/src/renderer/wg_engine/tvgWgRenderTarget.h b/src/renderer/wg_engine/tvgWgRenderTarget.h index 65dd6d4e..7d9ba590 100644 --- a/src/renderer/wg_engine/tvgWgRenderTarget.h +++ b/src/renderer/wg_engine/tvgWgRenderTarget.h @@ -52,10 +52,29 @@ public: void renderShape(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType); void renderPicture(WgContext& context, WgRenderDataPicture* renderData, WgPipelineBlendType blendType); + void renderClipPath(WgContext& context, WgRenderDataPaint* renderData); void clear(WGPUCommandEncoder commandEncoder); - void blend(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc, WgPipelineBlend* pipeline, WgBindGroupBlendMethod* blendMethod, WgBindGroupOpacity* opacity); - void composeBlend( + void blend( + WgContext& context, + WGPUCommandEncoder commandEncoder, + WgPipelineBlend* pipeline, + WgRenderStorage* targetSrc, + WgBindGroupBlendMethod* blendMethod, + WgBindGroupOpacity* opacity); + void blendMask( + WgContext& context, + WGPUCommandEncoder commandEncoder, + WgPipelineBlendMask* pipeline, + WgRenderStorage* texMsk, + WgRenderStorage* texSrc, + WgBindGroupBlendMethod* blendMethod, + WgBindGroupOpacity* opacity); + void maskCompose( + WgContext& context, + WGPUCommandEncoder commandEncoder, + WgRenderStorage* texMsk0); + void compose( WgContext& context, WGPUCommandEncoder commandEncoder, WgRenderStorage* texMsk, @@ -68,6 +87,9 @@ private: void drawShape(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType); void drawStroke(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType); + void drawShapeClipPath(WgContext& context, WgRenderDataShape* renderData); + void drawPictureClipPath(WgContext& context, WgRenderDataPicture* renderData); + void dispatchWorkgroups(WGPUComputePassEncoder computePassEncoder); WGPUComputePassEncoder beginComputePass(WGPUCommandEncoder commandEncoder); diff --git a/src/renderer/wg_engine/tvgWgRenderer.cpp b/src/renderer/wg_engine/tvgWgRenderer.cpp index e36e4a44..5ba138fa 100644 --- a/src/renderer/wg_engine/tvgWgRenderer.cpp +++ b/src/renderer/wg_engine/tvgWgRenderer.cpp @@ -67,6 +67,7 @@ void WgRenderer::release() mBlendMethodPool.release(mContext); mOpacityPool.release(mContext); mRenderStorageRoot.release(mContext); + mRenderStorageMask.release(mContext); mRenderStorageScreen.release(mContext); mRenderStorageInterm.release(mContext); mPipelines.release(); @@ -119,6 +120,9 @@ RenderData WgRenderer::prepare(const RenderShape& rshape, RenderData data, const if (rshape.stroke) renderDataShape->renderSettingsStroke.update(mContext, rshape.stroke->fill, rshape.stroke->color, flags); + // store clips data + renderDataShape->updateClips(clips); + return renderDataShape; } @@ -159,6 +163,9 @@ RenderData WgRenderer::prepare(Surface* surface, const RenderMesh* mesh, RenderD renderDataPicture->imageData.textureView); } + // store clips data + renderDataPicture->updateClips(clips); + return renderDataPicture; } @@ -175,6 +182,19 @@ bool WgRenderer::preRender() } +void WgRenderer::renderClipPath(Array& clips) +{ + for (uint32_t i = 0; i < clips.count; i++) { + renderClipPath(clips[i]->clips); + // render image to render target + mRenderStorageInterm.beginRenderPass(mCommandEncoder, true); + mRenderStorageInterm.renderClipPath(mContext, clips[i]); + mRenderStorageInterm.endRenderPass(); + mRenderStorageMask.maskCompose(mContext, mCommandEncoder, &mRenderStorageInterm); + } +} + + bool WgRenderer::renderShape(RenderData data) { // get current render storage @@ -183,11 +203,34 @@ bool WgRenderer::renderShape(RenderData data) WgPipelineBlendType blendType = WgPipelines::blendMethodToBlendType(mBlendMethod); WgRenderStorage* renderStorage = mRenderStorageStack.last(); assert(renderStorage); + // use masked blend + if (dataShape->clips.count > 0) { + // terminate current render pass + renderStorage->endRenderPass(); + // render clip path + mRenderStorageMask.beginRenderPass(mCommandEncoder, true); + mRenderStorageMask.renderClipPath(mContext, dataShape->clips[0]); + mRenderStorageMask.endRenderPass(); + renderClipPath(dataShape->clips); + // render image to render target + 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); + WgPipelineBlendMask* pipeline = &mContext.pipelines->computeBlendSolidMask; + if (dataShape->renderSettingsShape.fillType != WgRenderSettingsType::Solid) + pipeline = &mContext.pipelines->computeBlendGradientMask; + renderStorage->blendMask(mContext, mCommandEncoder, + pipeline, &mRenderStorageMask, &mRenderStorageInterm, blendMethod, opacity); + // restore current render pass + renderStorage->beginRenderPass(mCommandEncoder, false); // use hardware blend - if (WgPipelines::isBlendMethodSupportsHW(mBlendMethod)) + } else if (WgPipelines::isBlendMethodSupportsHW(mBlendMethod)) { renderStorage->renderShape(mContext, dataShape, blendType); // use custom blend - else { + } else { // terminate current render pass renderStorage->endRenderPass(); // render image to render target @@ -200,7 +243,8 @@ bool WgRenderer::renderShape(RenderData data) WgPipelineBlend* pipeline = &mContext.pipelines->computeBlendSolid; if (dataShape->renderSettingsShape.fillType != WgRenderSettingsType::Solid) pipeline = &mContext.pipelines->computeBlendGradient; - renderStorage->blend(mCommandEncoder, &mRenderStorageInterm, pipeline, blendMethod, opacity); + renderStorage->blend(mContext, mCommandEncoder, + pipeline, &mRenderStorageInterm, blendMethod, opacity); // restore current render pass renderStorage->beginRenderPass(mCommandEncoder, false); } @@ -210,26 +254,50 @@ bool WgRenderer::renderShape(RenderData data) bool WgRenderer::renderImage(RenderData data) { + // get current render storage + WgRenderDataPicture *dataPicture = (WgRenderDataPicture*)data; // get current render storage WgPipelineBlendType blendType = WgPipelines::blendMethodToBlendType(mBlendMethod); WgRenderStorage* renderStorage = mRenderStorageStack.last(); assert(renderStorage); + // use masked blend + if (dataPicture->clips.count > 0) { + // terminate current render pass + renderStorage->endRenderPass(); + // render clip path + mRenderStorageMask.beginRenderPass(mCommandEncoder, true); + mRenderStorageMask.renderClipPath(mContext, dataPicture->clips[0]); + mRenderStorageMask.endRenderPass(); + renderClipPath(dataPicture->clips); + // render image to render target + mRenderStorageInterm.beginRenderPass(mCommandEncoder, true); + mRenderStorageInterm.renderPicture(mContext, dataPicture, blendType); + mRenderStorageInterm.endRenderPass(); + // blend shape with current render storage + WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod); + WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, 255); + WgPipelineBlendMask* pipeline = &mContext.pipelines->computeBlendImageMask; + renderStorage->blendMask(mContext, mCommandEncoder, + pipeline, &mRenderStorageMask, &mRenderStorageInterm, blendMethod, opacity); + // restore current render pass + renderStorage->beginRenderPass(mCommandEncoder, false); // use hardware blend - if (WgPipelines::isBlendMethodSupportsHW(mBlendMethod)) - renderStorage->renderPicture(mContext, (WgRenderDataPicture *)data, blendType); + } else if (WgPipelines::isBlendMethodSupportsHW(mBlendMethod)) + renderStorage->renderPicture(mContext, dataPicture, blendType); // use custom blend else { // terminate current render pass renderStorage->endRenderPass(); // render image to render target mRenderStorageInterm.beginRenderPass(mCommandEncoder, true); - mRenderStorageInterm.renderPicture(mContext, (WgRenderDataPicture *)data, blendType); + mRenderStorageInterm.renderPicture(mContext, dataPicture, blendType); mRenderStorageInterm.endRenderPass(); // blend shape with current render storage WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod); WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, 255); WgPipelineBlend* pipeline = &mContext.pipelines->computeBlendImage; - renderStorage->blend(mCommandEncoder, &mRenderStorageInterm, pipeline, blendMethod, opacity); + renderStorage->blend(mContext, mCommandEncoder, + pipeline, &mRenderStorageInterm, blendMethod, opacity); // restore current render pass renderStorage->beginRenderPass(mCommandEncoder, false); } @@ -258,7 +326,7 @@ void WgRenderer::dispose(RenderData data) { RenderRegion WgRenderer::region(TVG_UNUSED RenderData data) { - return { 0, 0, INT32_MAX, INT32_MAX }; + return { 0, 0, (int32_t)mTargetSurface.w, (int32_t)mTargetSurface.h }; } @@ -355,6 +423,7 @@ bool WgRenderer::target(WGPUInstance instance, WGPUSurface surface, uint32_t w, initialize(); mRenderStorageInterm.initialize(mContext, w, h, WG_SSAA_SAMPLES); + mRenderStorageMask.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; @@ -400,7 +469,9 @@ bool WgRenderer::endComposite(TVG_UNUSED Compositor* cmp) // 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, &mContext.pipelines->computeBlendImage, blendMethod, opacity); + mRenderStorageStack.last()->blend(mContext, mCommandEncoder, + &mContext.pipelines->computeBlendImage, + renderStorageSrc, blendMethod, opacity); // back render targets to the pool mRenderStoragePool.free(mContext, renderStorageSrc); @@ -424,7 +495,7 @@ bool WgRenderer::endComposite(TVG_UNUSED Compositor* cmp) // compose and blend // dest = blend(dest, compose(src, msk, composeMethod), blendMethod, opacity) - renderStorageDst->composeBlend(mContext, mCommandEncoder, + renderStorageDst->compose(mContext, mCommandEncoder, renderStorageSrc, renderStorageMsk, composeMethod, blendMethod, opacity); diff --git a/src/renderer/wg_engine/tvgWgRenderer.h b/src/renderer/wg_engine/tvgWgRenderer.h index 4784a816..2553cb6b 100644 --- a/src/renderer/wg_engine/tvgWgRenderer.h +++ b/src/renderer/wg_engine/tvgWgRenderer.h @@ -62,12 +62,14 @@ private: void initialize(); void release(); void clearDisposes(); + void renderClipPath(Array& clips); - WGPUCommandEncoder mCommandEncoder{}; // render handles - WgRenderStorage mRenderStorageInterm; // intermidiate buffer to render - WgRenderStorage mRenderStorageRoot; // root render storage - WgRenderStorage mRenderStorageScreen; // storage with data after antializing - WgRenderStoragePool mRenderStoragePool; // pool to hold render tree storages + WGPUCommandEncoder mCommandEncoder{}; + WgRenderStorage mRenderStorageInterm; // intermidiate buffer to render + WgRenderStorage mRenderStorageMask; // buffer to render mask + WgRenderStorage mRenderStorageRoot; // root render storage + WgRenderStorage mRenderStorageScreen; // storage with data after antializing + WgRenderStoragePool mRenderStoragePool; // pool to hold render tree storages WgBindGroupOpacityPool mOpacityPool; // opacity, blend methods and composite methods pool WgBindGroupBlendMethodPool mBlendMethodPool; WgBindGroupCompositeMethodPool mCompositeMethodPool; diff --git a/src/renderer/wg_engine/tvgWgShaderSrc.cpp b/src/renderer/wg_engine/tvgWgShaderSrc.cpp index dfafe1f3..b1f40264 100644 --- a/src/renderer/wg_engine/tvgWgShaderSrc.cpp +++ b/src/renderer/wg_engine/tvgWgShaderSrc.cpp @@ -53,8 +53,8 @@ fn vs_main(in: VertexInput) -> VertexOutput { } @fragment -fn fs_main(in: VertexOutput) -> void { - // nothing to draw, just stencil value +fn fs_main(in: VertexOutput) -> @location(0) vec4f { + return vec4f(0.0, 0.0, 0.0, 1.0); } ); @@ -393,6 +393,34 @@ fn cs_main( @builtin(global_invocation_id) id: vec3u) { var Ra: f32 = 1.0; ); +const std::string strBlendMaskShaderHeader = 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; +@group(1) @binding(0) var blendMethod : u32; +@group(2) @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 colorMsk = textureLoad(imageMsk, id.xy); + if (colorMsk.a == 0.0) { return; } + let colorSrc = textureLoad(imageSrc, id.xy); + if (colorSrc.a == 0.0) { return; } + let colorDst = textureLoad(imageDst, id.xy); + + 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; +); + const std::string strBlendShaderPreConditionsGradient = WG_SHADER_SOURCE( Sc = Sc + Dc.rgb * (1.0 - Sa); Sa = Sa + Da * (1.0 - Sa); @@ -484,6 +512,31 @@ const std::string strComputeBlendImage = strBlendShaderFooter; const char* cShaderSource_PipelineComputeBlendImage = strComputeBlendImage.c_str(); +// pipeline shader modules blend solid mask +const std::string strComputeBlendSolidMask = + strBlendMaskShaderHeader + + strBlendShaderBlendMethod + + strBlendShaderFooter; +const char* cShaderSource_PipelineComputeBlendSolidMask = strComputeBlendSolidMask.c_str(); + +// pipeline shader modules blend gradient mask +const std::string strComputeBlendGradientMask = + strBlendMaskShaderHeader + + strBlendShaderPreConditionsGradient + + strBlendShaderBlendMethod + + strBlendShaderPostConditionsGradient + + strBlendShaderFooter; +const char* cShaderSource_PipelineComputeBlendGradientMask = strComputeBlendGradientMask.c_str(); + +// pipeline shader modules blend image mask +const std::string strComputeBlendImageMask = + strBlendMaskShaderHeader + + strBlendShaderPreConditionsImage + + strBlendShaderBlendMethod + + strBlendShaderPostConditionsImage + + strBlendShaderFooter; +const char* cShaderSource_PipelineComputeBlendImageMask = strComputeBlendImageMask.c_str(); + // pipeline shader modules clear const char* cShaderSource_PipelineComputeClear = WG_SHADER_SOURCE( @group(0) @binding(0) var imageDst : texture_storage_2d; @@ -494,8 +547,21 @@ fn cs_main( @builtin(global_invocation_id) id: vec3u) { } ); -// pipeline shader modules compose blend -const char* cShaderSource_PipelineComputeComposeBlend = WG_SHADER_SOURCE( +// pipeline shader modules compose +const char* cShaderSource_PipelineComputeMaskCompose = WG_SHADER_SOURCE( +@group(0) @binding(0) var imageMsk0 : texture_storage_2d; +@group(0) @binding(1) var imageMsk1 : texture_storage_2d; + +@compute @workgroup_size(8, 8) +fn cs_main( @builtin(global_invocation_id) id: vec3u) { + let colorMsk0 = textureLoad(imageMsk0, id.xy); + let colorMsk1 = textureLoad(imageMsk1, id.xy); + textureStore(imageMsk1, id.xy, colorMsk0 * colorMsk1); +} +); + +// pipeline shader modules compose +const char* cShaderSource_PipelineComputeCompose = 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; diff --git a/src/renderer/wg_engine/tvgWgShaderSrc.h b/src/renderer/wg_engine/tvgWgShaderSrc.h index 72243c0f..eb5f97bb 100644 --- a/src/renderer/wg_engine/tvgWgShaderSrc.h +++ b/src/renderer/wg_engine/tvgWgShaderSrc.h @@ -45,7 +45,11 @@ extern const char* cShaderSource_PipelineComputeClear; extern const char* cShaderSource_PipelineComputeBlendSolid; extern const char* cShaderSource_PipelineComputeBlendGradient; extern const char* cShaderSource_PipelineComputeBlendImage; -extern const char* cShaderSource_PipelineComputeComposeBlend; +extern const char* cShaderSource_PipelineComputeBlendSolidMask; +extern const char* cShaderSource_PipelineComputeBlendGradientMask; +extern const char* cShaderSource_PipelineComputeBlendImageMask; +extern const char* cShaderSource_PipelineComputeMaskCompose; +extern const char* cShaderSource_PipelineComputeCompose; extern const char* cShaderSource_PipelineComputeAntiAlias; #endif // _TVG_WG_SHADER_SRC_H_