From 0787e4663589f135c1dc17d2cdefb73c4940b63d Mon Sep 17 00:00:00 2001 From: Sergii Liebodkin Date: Wed, 14 Aug 2024 22:45:56 +0000 Subject: [PATCH] wg_engine: aabb based masking optimization * used fragment shaders for mask applience instead of compute shaders * less render targets swithing * shape aabb based on transformed shape bbox --- src/renderer/wg_engine/tvgWgCompositor.cpp | 81 +++++--- src/renderer/wg_engine/tvgWgCompositor.h | 9 +- src/renderer/wg_engine/tvgWgGeometry.h | 2 +- src/renderer/wg_engine/tvgWgPipelines.cpp | 121 ++++++----- src/renderer/wg_engine/tvgWgPipelines.h | 9 +- src/renderer/wg_engine/tvgWgRenderData.cpp | 14 ++ src/renderer/wg_engine/tvgWgRenderData.h | 2 + src/renderer/wg_engine/tvgWgRenderer.cpp | 15 +- src/renderer/wg_engine/tvgWgShaderSrc.cpp | 231 ++++++++++----------- src/renderer/wg_engine/tvgWgShaderSrc.h | 2 +- 10 files changed, 267 insertions(+), 219 deletions(-) diff --git a/src/renderer/wg_engine/tvgWgCompositor.cpp b/src/renderer/wg_engine/tvgWgCompositor.cpp index 7b639ee3..3b17467b 100755 --- a/src/renderer/wg_engine/tvgWgCompositor.cpp +++ b/src/renderer/wg_engine/tvgWgCompositor.cpp @@ -47,11 +47,17 @@ void WgCompositor::initialize(WgContext& context, uint32_t width, uint32_t heigh storageClipPath.initialize(context, width, height); storageInterm.initialize(context, width, height); storageDstCopy.initialize(context, width, height); + // composition and blend geometries + WgGeometryData geometryData; + geometryData.appendBlitBox(); + meshData.update(context, &geometryData); } void WgCompositor::release(WgContext& context) { + // composition and blend geometries + meshData.release(context); // release intermediate render storages storageInterm.release(context); storageDstCopy.release(context); @@ -252,6 +258,7 @@ void WgCompositor::blendImage(WgContext& context, WgRenderDataPicture* renderDat }; +// TODO: use direct mask applience void WgCompositor::composeShape(WgContext& context, WgRenderDataShape* renderData, WgRenderStorage* mask, CompositeMethod composeMethod) { assert(mask); @@ -268,16 +275,17 @@ void WgCompositor::composeShape(WgContext& context, WgRenderDataShape* renderDat beginRenderPass(commandEncoder, &storageInterm, true); drawShape(context, renderData, WgPipelineBlendType::Custom); endRenderPass(); - // run compisition - compose(commandEncoder, &storageInterm, mask, target, composeMethod); // restore current render pass beginRenderPass(commandEncoder, target, false); + RenderRegion rect { 0, 0,(int32_t)width, (int32_t)height }; + composeRegion(context, &storageInterm, mask, composeMethod, rect); } -void WgCompositor::composeStrokes(WgContext& context, WgRenderDataShape* renderData, WgRenderStorage* msk, CompositeMethod composeMethod) +// TODO: use direct mask applience +void WgCompositor::composeStrokes(WgContext& context, WgRenderDataShape* renderData, WgRenderStorage* mask, CompositeMethod composeMethod) { - assert(msk); + assert(mask); assert(renderData); assert(commandEncoder); assert(currentTarget); @@ -291,16 +299,17 @@ void WgCompositor::composeStrokes(WgContext& context, WgRenderDataShape* renderD beginRenderPass(commandEncoder, &storageInterm, true); drawStrokes(context, renderData, WgPipelineBlendType::Custom); endRenderPass(); - // run compisition - compose(commandEncoder, &storageInterm, msk, target, composeMethod); // restore current render pass beginRenderPass(commandEncoder, target, false); + RenderRegion rect { 0, 0, (int32_t)width, (int32_t)height }; + composeRegion(context, &storageInterm, mask, composeMethod, rect); } -void WgCompositor::composeImage(WgContext& context, WgRenderDataPicture* renderData, WgRenderStorage* msk, CompositeMethod composeMethod) +// TODO: use direct mask applience +void WgCompositor::composeImage(WgContext& context, WgRenderDataPicture* renderData, WgRenderStorage* mask, CompositeMethod composeMethod) { - assert(msk); + assert(mask); assert(renderData); assert(commandEncoder); assert(currentTarget); @@ -311,10 +320,38 @@ void WgCompositor::composeImage(WgContext& context, WgRenderDataPicture* renderD beginRenderPass(commandEncoder, &storageInterm, true); drawImage(context, renderData, WgPipelineBlendType::Custom); endRenderPass(); - // run compisition - compose(commandEncoder, &storageInterm, msk, target, composeMethod); // restore current render pass beginRenderPass(commandEncoder, target, false); + RenderRegion rect { 0, 0, (int32_t)width, (int32_t)height }; + composeRegion(context, &storageInterm, mask, composeMethod, rect); +} + + +void WgCompositor::composeScene(WgContext& context, WgRenderStorage* src, WgRenderStorage* mask, WgCompose* cmp) +{ + assert(cmp); + assert(src); + assert(mask); + assert(renderPassEncoder); + // cut viewport to screen dimensions + int32_t xmin = std::max(0, std::min((int32_t)width, cmp->aabb.x)); + int32_t ymin = std::max(0, std::min((int32_t)height, cmp->aabb.y)); + int32_t xmax = std::max(0, std::min((int32_t)width, cmp->aabb.x + cmp->aabb.w)); + int32_t ymax = std::max(0, std::min((int32_t)height, cmp->aabb.y + cmp->aabb.h)); + if ((xmin >= xmax) || (ymin >= ymax)) return; + RenderRegion rect { xmin, ymin, xmax - xmin, ymax - ymin }; + composeRegion(context, src, mask, cmp->method, rect); +} + + +void WgCompositor::composeRegion(WgContext& context, WgRenderStorage* src, WgRenderStorage* mask, CompositeMethod composeMethod, RenderRegion& rect) +{ + wgpuRenderPassEncoderSetScissorRect(renderPassEncoder, rect.x, rect.y, rect.w, rect.h); + wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 0); + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 0, src->bindGroupTexure, 0, nullptr); + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 1, mask->bindGroupTexure, 0, nullptr); + wgpuRenderPassEncoderSetPipeline(renderPassEncoder, pipelines->sceneComp[(uint32_t)composeMethod]); + meshData.drawImage(context, renderPassEncoder); } @@ -483,27 +520,3 @@ void WgCompositor::blend(WGPUCommandEncoder encoder, WgRenderStorage* src, WgRen wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, (width + 7) / 8, (height + 7) / 8, 1); wgpuComputePassEncoderEnd(computePassEncoder); } - - -void WgCompositor::compose(WGPUCommandEncoder encoder, WgRenderStorage* src, WgRenderStorage* mask, WgRenderStorage* dst, CompositeMethod composeMethod) -{ - assert(src); - assert(mask); - assert(dst); - assert(!renderPassEncoder); - // copy dst storage to temporary read only storage - const WGPUImageCopyTexture texSrc { .texture = dst->texture }; - const WGPUImageCopyTexture texDst { .texture = storageDstCopy.texture }; - const WGPUExtent3D copySize { .width = width, .height = height, .depthOrArrayLayers = 1 }; - wgpuCommandEncoderCopyTextureToTexture(encoder, &texSrc, &texDst, ©Size); - // execute compose shader - const WGPUComputePassDescriptor computePassDescriptor{}; - WGPUComputePassEncoder computePassEncoder = wgpuCommandEncoderBeginComputePass(encoder, &computePassDescriptor); - wgpuComputePassEncoderSetBindGroup(computePassEncoder, 0, src->bindGroupRead, 0, nullptr); - wgpuComputePassEncoderSetBindGroup(computePassEncoder, 1, mask->bindGroupRead, 0, nullptr); - wgpuComputePassEncoderSetBindGroup(computePassEncoder, 2, storageDstCopy.bindGroupRead, 0, nullptr); - wgpuComputePassEncoderSetBindGroup(computePassEncoder, 3, dst->bindGroupWrite, 0, nullptr); - wgpuComputePassEncoderSetPipeline(computePassEncoder, pipelines->compose[(size_t)composeMethod]); - wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, (width + 7) / 8, (height + 7) / 8, 1); - wgpuComputePassEncoderEnd(computePassEncoder); -} diff --git a/src/renderer/wg_engine/tvgWgCompositor.h b/src/renderer/wg_engine/tvgWgCompositor.h index b4d6a409..5905c36c 100755 --- a/src/renderer/wg_engine/tvgWgCompositor.h +++ b/src/renderer/wg_engine/tvgWgCompositor.h @@ -27,7 +27,8 @@ #include "tvgWgRenderData.h" struct WgCompose: public Compositor { - BlendMethod blend; + BlendMethod blend{}; + RenderRegion aabb{}; }; class WgCompositor { @@ -51,8 +52,12 @@ private: WgRenderStorage storageInterm; WgRenderStorage storageClipPath; WgRenderStorage storageDstCopy; + // composition and blend geometries + WgMeshData meshData; static WgPipelineBlendType blendMethodToBlendType(BlendMethod blendMethod); + + void composeRegion(WgContext& context, WgRenderStorage* src, WgRenderStorage* mask, CompositeMethod composeMethod, RenderRegion& rect); public: // render target dimensions uint32_t width{}; @@ -76,6 +81,7 @@ public: void composeShape(WgContext& context, WgRenderDataShape* renderData, WgRenderStorage* mask, CompositeMethod composeMethod); void composeStrokes(WgContext& context, WgRenderDataShape* renderData, WgRenderStorage* mask, CompositeMethod composeMethod); void composeImage(WgContext& context, WgRenderDataPicture* renderData, WgRenderStorage* mask, CompositeMethod composeMethod); + void composeScene(WgContext& context, WgRenderStorage* src, WgRenderStorage* mask, WgCompose* cmp); void drawClipPath(WgContext& context, WgRenderDataShape* renderData); void drawShape(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType); @@ -84,7 +90,6 @@ public: void mergeMasks(WGPUCommandEncoder encoder, WgRenderStorage* mask0, WgRenderStorage* mask1); void blend(WGPUCommandEncoder encoder, WgRenderStorage* src, WgRenderStorage* dst, uint8_t opacity, BlendMethod blendMethod, WgRenderRasterType rasterType); - void compose(WGPUCommandEncoder encoder, WgRenderStorage* src, WgRenderStorage* mask, WgRenderStorage* dst, CompositeMethod composeMethod); }; #endif // _TVG_WG_COMPOSITOR_H_ diff --git a/src/renderer/wg_engine/tvgWgGeometry.h b/src/renderer/wg_engine/tvgWgGeometry.h index c9d4bf2b..630b0932 100755 --- a/src/renderer/wg_engine/tvgWgGeometry.h +++ b/src/renderer/wg_engine/tvgWgGeometry.h @@ -59,7 +59,7 @@ public: inline void normalize() { float rlen = 1.0f / length(); x *= rlen; y *= rlen; } inline WgPoint normal() const { float rlen = 1.0f / length(); return { x * rlen, y * rlen }; } inline WgPoint lerp(const WgPoint& p, float t) const { return { x + (p.x - x) * t, y + (p.y - y) * t }; }; - inline WgPoint trans(const Matrix& m) const { return { x * m.e11 + y * m.e12, x * m.e21 + y * m.e22 }; }; + inline WgPoint trans(const Matrix& m) const { return { x * m.e11 + y * m.e12 + m.e13, x * m.e21 + y * m.e22 + m.e23 }; }; }; struct WgMath diff --git a/src/renderer/wg_engine/tvgWgPipelines.cpp b/src/renderer/wg_engine/tvgWgPipelines.cpp index 465844bb..260a0889 100755 --- a/src/renderer/wg_engine/tvgWgPipelines.cpp +++ b/src/renderer/wg_engine/tvgWgPipelines.cpp @@ -43,7 +43,8 @@ WGPUPipelineLayout WgPipelines::createPipelineLayout(WGPUDevice device, const WG WGPURenderPipeline WgPipelines::createRenderPipeline( WGPUDevice device, const char* pipelineLabel, - const WGPUShaderModule shaderModule, const WGPUPipelineLayout pipelineLayout, + const WGPUShaderModule shaderModule, const char* vsEntryPoint, const char* fsEntryPoint, + const WGPUPipelineLayout pipelineLayout, const WGPUVertexBufferLayout *vertexBufferLayouts, const uint32_t vertexBufferLayoutsCount, const WGPUColorWriteMaskFlags writeMask, const WGPUCompareFunction stencilFunctionFrnt, const WGPUStencilOperation stencilOperationFrnt, @@ -58,8 +59,8 @@ WGPURenderPipeline WgPipelines::createRenderPipeline( .stencilBack = { .compare = stencilFunctionBack, .failOp = stencilOperationBack, .depthFailOp = stencilOperationBack, .passOp = stencilOperationBack }, .stencilReadMask = 0xFFFFFFFF, .stencilWriteMask = 0xFFFFFFFF }; - const WGPUVertexState vertexState { .module = shaderModule, .entryPoint = "vs_main", .bufferCount = vertexBufferLayoutsCount, .buffers = vertexBufferLayouts }; - const WGPUFragmentState fragmentState { .module = shaderModule, .entryPoint = "fs_main", .targetCount = 1, .targets = colorTargetStates }; + const WGPUVertexState vertexState { .module = shaderModule, .entryPoint = vsEntryPoint, .bufferCount = vertexBufferLayoutsCount, .buffers = vertexBufferLayouts }; + const WGPUFragmentState fragmentState { .module = shaderModule, .entryPoint = fsEntryPoint, .targetCount = 1, .targets = colorTargetStates }; const WGPURenderPipelineDescriptor renderPipelineDesc { .label = pipelineLabel, .layout = pipelineLayout, @@ -182,6 +183,10 @@ void WgPipelines::initialize(WgContext& context) layouts.layoutBuffer2Un, layouts.layoutTexSampled }; + const WGPUBindGroupLayout bindGroupLayoutsSceneComp[] = { + layouts.layoutTexSampled, + layouts.layoutTexSampled + }; const WGPUBindGroupLayout bindGroupLayoutsMergeMasks[] = { layouts.layoutTexStrorage1RO, layouts.layoutTexStrorage1RO, @@ -193,12 +198,6 @@ void WgPipelines::initialize(WgContext& context) layouts.layoutTexStrorage1WO, layouts.layoutBuffer1Un }; - const WGPUBindGroupLayout bindGroupLayoutsCompose[] = { - layouts.layoutTexStrorage1RO, - layouts.layoutTexStrorage1RO, - layouts.layoutTexStrorage1RO, - layouts.layoutTexStrorage1WO - }; const WGPUBindGroupLayout bindGroupLayoutsCopy[] = { layouts.layoutTexStrorage1RO, layouts.layoutTexScreen1WO @@ -209,8 +208,8 @@ void WgPipelines::initialize(WgContext& context) layoutSolid = createPipelineLayout(context.device, bindGroupLayoutsSolid, 3); layoutGradient = createPipelineLayout(context.device, bindGroupLayoutsGradient, 3); layoutImage = createPipelineLayout(context.device, bindGroupLayoutsImage, 3); + layoutSceneComp = createPipelineLayout(context.device, bindGroupLayoutsSceneComp, 2); layoutBlend = createPipelineLayout(context.device, bindGroupLayoutsBlend, 4); - layoutCompose = createPipelineLayout(context.device, bindGroupLayoutsCompose, 4); layoutMergeMasks = createPipelineLayout(context.device, bindGroupLayoutsMergeMasks, 3); layoutCopy = createPipelineLayout(context.device, bindGroupLayoutsCopy, 2); @@ -220,6 +219,7 @@ void WgPipelines::initialize(WgContext& context) shaderRadial = createShaderModule(context.device, "The shader radial", cShaderSrc_Radial); shaderLinear = createShaderModule(context.device, "The shader linear", cShaderSrc_Linear); shaderImage = createShaderModule(context.device, "The shader image", cShaderSrc_Image); + shaderSceneComp = createShaderModule(context.device, "The shader scene composition", cShaderSrc_Scene_Comp); // computes shader modules shaderMergeMasks = createShaderModule(context.device, "The shader merge mask", cShaderSrc_MergeMasks); shaderCopy = createShaderModule(context.device, "The shader copy", cShaderSrc_Copy); @@ -227,7 +227,7 @@ void WgPipelines::initialize(WgContext& context) // render pipeline winding winding = createRenderPipeline( context.device, "The render pipeline winding", - shaderStencil, layoutStencil, + shaderStencil, "vs_main", "fs_main", layoutStencil, vertexBufferLayoutsShape, 1, WGPUColorWriteMask_None, WGPUCompareFunction_Always, WGPUStencilOperation_IncrementWrap, @@ -236,7 +236,7 @@ void WgPipelines::initialize(WgContext& context) // render pipeline even-odd evenodd = createRenderPipeline( context.device, "The render pipeline even-odd", - shaderStencil, layoutStencil, + shaderStencil, "vs_main", "fs_main", layoutStencil, vertexBufferLayoutsShape, 1, WGPUColorWriteMask_None, WGPUCompareFunction_Always, WGPUStencilOperation_Invert, @@ -245,7 +245,7 @@ void WgPipelines::initialize(WgContext& context) // render pipeline direct direct = createRenderPipeline( context.device, "The render pipeline direct", - shaderStencil, layoutStencil, + shaderStencil,"vs_main", "fs_main", layoutStencil, vertexBufferLayoutsShape, 1, WGPUColorWriteMask_None, WGPUCompareFunction_Always, WGPUStencilOperation_Replace, @@ -254,7 +254,7 @@ void WgPipelines::initialize(WgContext& context) // render pipeline clip path clipPath = createRenderPipeline( context.device, "The render pipeline clip path", - shaderStencil, layoutStencil, + shaderStencil, "vs_main", "fs_main", layoutStencil, vertexBufferLayoutsShape, 1, WGPUColorWriteMask_All, WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero, @@ -265,7 +265,7 @@ void WgPipelines::initialize(WgContext& context) for (uint32_t i = 0; i < 3; i++) { solid[i] = createRenderPipeline( context.device, "The render pipeline solid", - shaderSolid, layoutSolid, + shaderSolid, "vs_main", "fs_main", layoutSolid, vertexBufferLayoutsShape, 1, WGPUColorWriteMask_All, WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero, @@ -277,7 +277,7 @@ void WgPipelines::initialize(WgContext& context) for (uint32_t i = 0; i < 3; i++) { radial[i] = createRenderPipeline( context.device, "The render pipeline radial", - shaderRadial, layoutGradient, + shaderRadial, "vs_main", "fs_main", layoutGradient, vertexBufferLayoutsShape, 1, WGPUColorWriteMask_All, WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero, @@ -289,7 +289,7 @@ void WgPipelines::initialize(WgContext& context) for (uint32_t i = 0; i < 3; i++) { linear[i] = createRenderPipeline( context.device, "The render pipeline linear", - shaderLinear, layoutGradient, + shaderLinear, "vs_main", "fs_main", layoutGradient, vertexBufferLayoutsShape, 1, WGPUColorWriteMask_All, WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero, @@ -301,7 +301,7 @@ void WgPipelines::initialize(WgContext& context) for (uint32_t i = 0; i < 3; i++) { image[i] = createRenderPipeline( context.device, "The render pipeline image", - shaderImage, layoutImage, + shaderImage, "vs_main", "fs_main", layoutImage, vertexBufferLayoutsImage, 2, WGPUColorWriteMask_All, WGPUCompareFunction_Always, WGPUStencilOperation_Zero, @@ -309,6 +309,55 @@ void WgPipelines::initialize(WgContext& context) primitiveState, multisampleState, blendStates[i]); } + // compute shader names + const char* shaderComposeNames[] { + "fs_main_None", + "fs_main_ClipPath", + "fs_main_AlphaMask", + "fs_main_InvAlphaMask", + "fs_main_LumaMask", + "fs_main_InvLumaMask", + "fs_main_AddMask", + "fs_main_SubtractMask", + "fs_main_IntersectMask", + "fs_main_DifferenceMask", + "fs_main_LightenMask", + "fs_main_DarkenMask" + }; + + // compose shader blend states + const WGPUBlendState composeBlends[] { + blendStateNrm, // None + blendStateNrm, // ClipPath + blendStateNrm, // AlphaMask + blendStateNrm, // InvAlphaMask + blendStateNrm, // LumaMask + blendStateNrm, // InvLumaMask + blendStateSrc, // AddMask + blendStateSrc, // SubtractMask + blendStateSrc, // IntersectMask + blendStateSrc, // DifferenceMask + blendStateSrc, // LightenMask + blendStateSrc // DarkenMask + }; + + // render pipeline scene composition + size_t shaderComposeNamesCnt = sizeof(shaderComposeNames) / sizeof(shaderComposeNames[0]); + size_t composeBlendspCnt = sizeof(composeBlends) / sizeof(composeBlends[0]); + size_t sceneCompCnt = sizeof(sceneComp) / sizeof(sceneComp[0]); + assert(shaderComposeNamesCnt == composeBlendspCnt); + assert(composeBlendspCnt == sceneCompCnt); + for (uint32_t i = 0; i < sceneCompCnt; i++) { + sceneComp[i] = createRenderPipeline( + context.device, "The render pipeline scene composition", + shaderSceneComp, "vs_main", shaderComposeNames[i], layoutSceneComp, + vertexBufferLayoutsImage, 2, + WGPUColorWriteMask_All, + WGPUCompareFunction_Always, WGPUStencilOperation_Zero, + WGPUCompareFunction_Always, WGPUStencilOperation_Zero, + primitiveState, multisampleState, composeBlends[i]); + } + // compute pipelines mergeMasks = createComputePipeline(context.device, "The pipeline merge masks", shaderMergeMasks, "cs_main", layoutMergeMasks); copy = createComputePipeline(context.device, "The pipeline copy", shaderCopy, "cs_main", layoutCopy); @@ -346,37 +395,14 @@ void WgPipelines::initialize(WgContext& context) blendGradient[i] = createComputePipeline(context.device, "The pipeline blend gradient", shaderBlendGradient, shaderBlendNames[i], layoutBlend); blendImage[i] = createComputePipeline(context.device, "The pipeline blend image", shaderBlendImage, shaderBlendNames[i], layoutBlend); } - - // compute shader compose names - const char* shaderComposeNames[] { - "cs_main_None", - "cs_main_ClipPath", - "cs_main_AlphaMask", - "cs_main_InvAlphaMask", - "cs_main_LumaMask", - "cs_main_InvLumaMask", - "cs_main_AddMask", - "cs_main_SubtractMask", - "cs_main_IntersectMask", - "cs_main_DifferenceMask", - "cs_main_LightenMask", - "cs_main_DarkenMask" - }; - - // create compose shaders - shaderCompose = createShaderModule(context.device, "The shader compose", cShaderSrc_Compose); - - // create compose pipelines - const size_t shaderComposeNamesCnt = sizeof(shaderComposeNames) / sizeof(shaderComposeNames[0]); - const size_t pipesComposeCnt = sizeof(compose) / sizeof(compose[0]); - assert(shaderComposeNamesCnt == pipesComposeCnt); - for (uint32_t i = 0; i < pipesComposeCnt; i++) - compose[i] = createComputePipeline(context.device, "The pipeline compose", shaderCompose, shaderComposeNames[i], layoutCompose); } void WgPipelines::releaseGraphicHandles(WgContext& context) { releaseRenderPipeline(clipPath); + size_t pipesSceneCompCnt = sizeof(sceneComp) / sizeof(sceneComp[0]); + for (uint32_t i = 0; i < pipesSceneCompCnt; i++) + releaseRenderPipeline(sceneComp[i]); for (uint32_t i = 0; i < 3; i++) { releaseRenderPipeline(image[i]); releaseRenderPipeline(linear[i]); @@ -386,10 +412,12 @@ void WgPipelines::releaseGraphicHandles(WgContext& context) releaseRenderPipeline(direct); releaseRenderPipeline(evenodd); releaseRenderPipeline(winding); + releasePipelineLayout(layoutSceneComp); releasePipelineLayout(layoutImage); releasePipelineLayout(layoutGradient); releasePipelineLayout(layoutSolid); releasePipelineLayout(layoutStencil); + releaseShaderModule(shaderSceneComp); releaseShaderModule(shaderImage); releaseShaderModule(shaderLinear); releaseShaderModule(shaderRadial); @@ -400,9 +428,6 @@ void WgPipelines::releaseGraphicHandles(WgContext& context) void WgPipelines::releaseComputeHandles(WgContext& context) { - size_t pipesComposeCnt = sizeof(compose) / sizeof(compose[0]); - for (uint32_t i = 0; i < pipesComposeCnt; i++) - releaseComputePipeline(compose[i]); const size_t pipesBlendCnt = sizeof(blendSolid)/sizeof(blendSolid[0]); for (uint32_t i = 0; i < pipesBlendCnt; i++) { releaseComputePipeline(blendImage[i]); @@ -411,11 +436,9 @@ void WgPipelines::releaseComputeHandles(WgContext& context) } releaseComputePipeline(copy); releaseComputePipeline(mergeMasks); - releasePipelineLayout(layoutCompose); releasePipelineLayout(layoutBlend); releasePipelineLayout(layoutMergeMasks); releasePipelineLayout(layoutCopy); - releaseShaderModule(shaderCompose); releaseShaderModule(shaderBlendImage); releaseShaderModule(shaderBlendGradient); releaseShaderModule(shaderBlendSolid); diff --git a/src/renderer/wg_engine/tvgWgPipelines.h b/src/renderer/wg_engine/tvgWgPipelines.h index 4cb84d6c..57320b45 100755 --- a/src/renderer/wg_engine/tvgWgPipelines.h +++ b/src/renderer/wg_engine/tvgWgPipelines.h @@ -35,12 +35,12 @@ private: WGPUShaderModule shaderRadial{}; WGPUShaderModule shaderLinear{}; WGPUShaderModule shaderImage{}; + WGPUShaderModule shaderSceneComp{}; // compute pipeline shaders WGPUShaderModule shaderMergeMasks; WGPUShaderModule shaderBlendSolid; WGPUShaderModule shaderBlendGradient; WGPUShaderModule shaderBlendImage; - WGPUShaderModule shaderCompose; WGPUShaderModule shaderCopy; private: // graphics pipeline layouts @@ -48,10 +48,10 @@ private: WGPUPipelineLayout layoutSolid{}; WGPUPipelineLayout layoutGradient{}; WGPUPipelineLayout layoutImage{}; + WGPUPipelineLayout layoutSceneComp{}; // compute pipeline layouts WGPUPipelineLayout layoutMergeMasks{}; WGPUPipelineLayout layoutBlend{}; - WGPUPipelineLayout layoutCompose{}; WGPUPipelineLayout layoutCopy{}; public: // graphics pipeline @@ -63,13 +63,13 @@ public: WGPURenderPipeline radial[3]{}; WGPURenderPipeline linear[3]{}; WGPURenderPipeline image[3]{}; + WGPURenderPipeline sceneComp[12]; WGPURenderPipeline clipPath{}; // compute pipeline WGPUComputePipeline mergeMasks; WGPUComputePipeline blendSolid[14]; WGPUComputePipeline blendGradient[14]; WGPUComputePipeline blendImage[14]; - WGPUComputePipeline compose[12]; WGPUComputePipeline copy; private: void releaseGraphicHandles(WgContext& context); @@ -79,7 +79,8 @@ private: WGPUPipelineLayout createPipelineLayout(WGPUDevice device, const WGPUBindGroupLayout* bindGroupLayouts, const uint32_t bindGroupLayoutsCount); WGPURenderPipeline createRenderPipeline( WGPUDevice device, const char* pipelineLabel, - const WGPUShaderModule shaderModule, const WGPUPipelineLayout pipelineLayout, + const WGPUShaderModule shaderModule, const char* vsEntryPoint, const char* fsEntryPoint, + const WGPUPipelineLayout pipelineLayout, const WGPUVertexBufferLayout *vertexBufferLayouts, const uint32_t vertexBufferLayoutsCount, const WGPUColorWriteMaskFlags writeMask, const WGPUCompareFunction stencilFunctionFrnt, const WGPUStencilOperation stencilOperationFrnt, diff --git a/src/renderer/wg_engine/tvgWgRenderData.cpp b/src/renderer/wg_engine/tvgWgRenderData.cpp index 56eabc08..16840b63 100755 --- a/src/renderer/wg_engine/tvgWgRenderData.cpp +++ b/src/renderer/wg_engine/tvgWgRenderData.cpp @@ -306,6 +306,18 @@ void WgRenderDataShape::updateBBox(WgPoint pmin, WgPoint pmax) } +void WgRenderDataShape::updateAABB(const Matrix& rt) { + WgPoint p0 = WgPoint(pMin.x, pMin.y).trans(rt); + WgPoint p1 = WgPoint(pMax.x, pMin.y).trans(rt); + WgPoint p2 = WgPoint(pMin.x, pMax.y).trans(rt); + WgPoint p3 = WgPoint(pMax.x, pMax.y).trans(rt); + aabb.x = std::min({p0.x, p1.x, p2.x, p3.x}); + aabb.y = std::min({p0.y, p1.y, p2.y, p3.y}); + aabb.w = std::max({p0.x, p1.x, p2.x, p3.x}) - aabb.x; + aabb.h = std::max({p0.y, p1.y, p2.y, p3.y}) - aabb.y; +} + + void WgRenderDataShape::updateMeshes(WgContext &context, const RenderShape &rshape, const Matrix& rt) { releaseMeshes(context); @@ -368,6 +380,7 @@ void WgRenderDataShape::updateMeshes(WgContext &context, const RenderShape &rsha for (uint32_t i = 0; i < polylines.count; i++) delete polylines[i]; // update shapes bbox + updateAABB(rt); meshDataBBox.update(context, pMin, pMax); } @@ -443,6 +456,7 @@ void WgRenderDataShape::updateStrokes(WgContext& context, const WgPolyline* poly geometryData.positions.getBBox(pmin, pmax); meshGroupStrokes.append(context, &geometryData); meshGroupStrokesBBox.append(context, pmin, pmax); + updateBBox(pmin, pmax); } } } diff --git a/src/renderer/wg_engine/tvgWgRenderData.h b/src/renderer/wg_engine/tvgWgRenderData.h index f0837be9..ed8afc1b 100755 --- a/src/renderer/wg_engine/tvgWgRenderData.h +++ b/src/renderer/wg_engine/tvgWgRenderData.h @@ -97,6 +97,7 @@ struct WgRenderDataPaint WGPUBuffer bufferBlendSettings{}; WGPUBindGroup bindGroupPaint{}; RenderRegion viewport{}; + RenderRegion aabb{}; float opacity{}; Array clips; @@ -123,6 +124,7 @@ struct WgRenderDataShape: public WgRenderDataPaint FillRule fillRule{}; void updateBBox(WgPoint pmin, WgPoint pmax); + void updateAABB(const Matrix& rt); void updateMeshes(WgContext& context, const RenderShape& rshape, const Matrix& rt); void updateShapes(WgContext& context, const WgPolyline* polyline); void updateStrokesList(WgContext& context, Array polylines, const RenderStroke* rstroke, float totalLen, float trimBegin, float trimEnd); diff --git a/src/renderer/wg_engine/tvgWgRenderer.cpp b/src/renderer/wg_engine/tvgWgRenderer.cpp index 5ce90ef6..7c0fab52 100755 --- a/src/renderer/wg_engine/tvgWgRenderer.cpp +++ b/src/renderer/wg_engine/tvgWgRenderer.cpp @@ -199,8 +199,12 @@ void WgRenderer::dispose(RenderData data) { } -RenderRegion WgRenderer::region(TVG_UNUSED RenderData data) +RenderRegion WgRenderer::region(RenderData data) { + auto renderData = (WgRenderDataPaint*)data; + if (renderData->type() == Type::Shape) { + return renderData->aabb; + } return { 0, 0, (int32_t)mTargetSurface.w, (int32_t)mTargetSurface.h }; } @@ -316,6 +320,7 @@ bool WgRenderer::target(WGPUInstance instance, WGPUSurface surface, uint32_t w, Compositor* WgRenderer::target(TVG_UNUSED const RenderRegion& region, TVG_UNUSED ColorSpace cs) { mCompositorStack.push(new WgCompose); + mCompositorStack.last()->aabb = region; return mCompositorStack.last(); } @@ -354,6 +359,8 @@ bool WgRenderer::endComposite(Compositor* cmp) mCompositor.blend(mCommandEncoder, src, dst, comp->opacity, comp->blend, WgRenderRasterType::Image); // back render targets to the pool mRenderStoragePool.free(mContext, src); + // begin previous render pass + mCompositor.beginRenderPass(mCommandEncoder, mRenderStorageStack.last(), false); } else { // finish composition // get source, mask and destination render storages WgRenderStorage* src = mRenderStorageStack.last(); @@ -361,14 +368,14 @@ bool WgRenderer::endComposite(Compositor* cmp) WgRenderStorage* msk = mRenderStorageStack.last(); mRenderStorageStack.pop(); WgRenderStorage* dst = mRenderStorageStack.last(); + // begin previous render pass + mCompositor.beginRenderPass(mCommandEncoder, dst, false); // apply composition - mCompositor.compose(mCommandEncoder, src, msk, dst,comp->method); + mCompositor.composeScene(mContext, src, msk, comp); // back render targets to the pool mRenderStoragePool.free(mContext, src); mRenderStoragePool.free(mContext, msk); } - // begin previous render pass - mCompositor.beginRenderPass(mCommandEncoder, mRenderStorageStack.last(), false); // delete current compositor settings delete mCompositorStack.last(); diff --git a/src/renderer/wg_engine/tvgWgShaderSrc.cpp b/src/renderer/wg_engine/tvgWgShaderSrc.cpp index 5e2aaa5c..bb191720 100755 --- a/src/renderer/wg_engine/tvgWgShaderSrc.cpp +++ b/src/renderer/wg_engine/tvgWgShaderSrc.cpp @@ -184,6 +184,113 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f { }; )"; +//************************************************************************ +// graphics shader source: scene compose +//************************************************************************ + +const char* cShaderSrc_Scene_Comp = R"( +struct VertexInput { @location(0) position: vec2f, @location(1) texCoord: vec2f }; +struct VertexOutput { @builtin(position) position: vec4f, @location(0) texCoord: vec2f }; + +@group(0) @binding(0) var uSamplerSrc : sampler; +@group(0) @binding(1) var uTextureSrc : texture_2d; +@group(1) @binding(0) var uSamplerMsk : sampler; +@group(1) @binding(1) var uTextureMsk : texture_2d; + +@vertex +fn vs_main(in: VertexInput) -> VertexOutput { + var out: VertexOutput; + out.position = vec4f(in.position.xy, 0.0, 1.0); + out.texCoord = in.texCoord; + return out; +} + +@fragment +fn fs_main_None(in: VertexOutput) -> @location(0) vec4f { + let src: vec4f = textureSample(uTextureSrc, uSamplerSrc, in.texCoord.xy); + return vec4f(src); +}; + +@fragment +fn fs_main_ClipPath(in: VertexOutput) -> @location(0) vec4f { + let src: vec4f = textureSample(uTextureSrc, uSamplerSrc, in.texCoord.xy); + let msk: vec4f = textureSample(uTextureMsk, uSamplerMsk, in.texCoord.xy); + return vec4f(src * msk.a); +}; + +@fragment +fn fs_main_AlphaMask(in: VertexOutput) -> @location(0) vec4f { + let src: vec4f = textureSample(uTextureSrc, uSamplerSrc, in.texCoord.xy); + let msk: vec4f = textureSample(uTextureMsk, uSamplerMsk, in.texCoord.xy); + return vec4f(src * msk.a); +}; + +@fragment +fn fs_main_InvAlphaMask(in: VertexOutput) -> @location(0) vec4f { + let src: vec4f = textureSample(uTextureSrc, uSamplerSrc, in.texCoord.xy); + let msk: vec4f = textureSample(uTextureMsk, uSamplerMsk, in.texCoord.xy); + return vec4f(src * (1.0 - msk.a)); +}; + +@fragment +fn fs_main_LumaMask(in: VertexOutput) -> @location(0) vec4f { + let src: vec4f = textureSample(uTextureSrc, uSamplerSrc, in.texCoord.xy); + let msk: vec4f = textureSample(uTextureMsk, uSamplerMsk, in.texCoord.xy); + let luma: f32 = dot(msk.rgb, vec3f(0.2125, 0.7154, 0.0721)); + return vec4f(src * luma); +}; + +@fragment +fn fs_main_InvLumaMask(in: VertexOutput) -> @location(0) vec4f { + let src: vec4f = textureSample(uTextureSrc, uSamplerSrc, in.texCoord.xy); + let msk: vec4f = textureSample(uTextureMsk, uSamplerMsk, in.texCoord.xy); + let luma: f32 = dot(msk.rgb, vec3f(0.2125, 0.7154, 0.0721)); + return vec4f(src * (1.0 - luma)); +}; + +@fragment +fn fs_main_AddMask(in: VertexOutput) -> @location(0) vec4f { + let src: vec4f = textureSample(uTextureSrc, uSamplerSrc, in.texCoord.xy); + let msk: vec4f = textureSample(uTextureMsk, uSamplerMsk, in.texCoord.xy); + return vec4f(src.rgb, src.a + msk.a * (1.0 - src.a)); +}; + +@fragment +fn fs_main_SubtractMask(in: VertexOutput) -> @location(0) vec4f { + let src: vec4f = textureSample(uTextureSrc, uSamplerSrc, in.texCoord.xy); + let msk: vec4f = textureSample(uTextureMsk, uSamplerMsk, in.texCoord.xy); + return vec4f(src.rgb, src.a * (1.0 - msk.a)); +}; + +@fragment +fn fs_main_IntersectMask(in: VertexOutput) -> @location(0) vec4f { + let src: vec4f = textureSample(uTextureSrc, uSamplerSrc, in.texCoord.xy); + let msk: vec4f = textureSample(uTextureMsk, uSamplerMsk, in.texCoord.xy); + return vec4f(src.rgb, src.a * msk.a); +}; + +@fragment +fn fs_main_DifferenceMask(in: VertexOutput) -> @location(0) vec4f { + let src: vec4f = textureSample(uTextureSrc, uSamplerSrc, in.texCoord.xy); + let msk: vec4f = textureSample(uTextureMsk, uSamplerMsk, in.texCoord.xy); + return vec4f(src.rgb, src.a * (1.0 - msk.a) + msk.a * (1.0 - src.a)); +}; + +@fragment +fn fs_main_LightenMask(in: VertexOutput) -> @location(0) vec4f { + let src: vec4f = textureSample(uTextureSrc, uSamplerSrc, in.texCoord.xy); + let msk: vec4f = textureSample(uTextureMsk, uSamplerMsk, in.texCoord.xy); + return vec4f(src.rgb, max(src.a, msk.a)); +}; + +@fragment +fn fs_main_DarkenMask(in: VertexOutput) -> @location(0) vec4f { + let src: vec4f = textureSample(uTextureSrc, uSamplerSrc, in.texCoord.xy); + let msk: vec4f = textureSample(uTextureMsk, uSamplerMsk, in.texCoord.xy); + return vec4f(src.rgb, min(src.a, msk.a)); +}; +)"; + //************************************************************************ // compute shader source: merge clip path masks //************************************************************************ @@ -412,130 +519,6 @@ fn cs_main_SoftLight(@builtin(global_invocation_id) id: vec3u) { }; )"; -//************************************************************************ -// compute shader source: compose -//************************************************************************ - -const char* cShaderSrc_Compose = R"( -@group(0) @binding(0) var imageSrc : texture_storage_2d; -@group(1) @binding(0) var imageMsk : texture_storage_2d; -@group(2) @binding(0) var imageDst : texture_storage_2d; -@group(3) @binding(0) var imageTgt : texture_storage_2d; - -struct FragData { Sc: vec3f, Sa: f32, Mc: vec3f, Ma: f32, Dc: vec3f, Da: f32 }; -fn getFragData(id: vec2u) -> FragData { - let colorSrc = textureLoad(imageSrc, id.xy); - let colorMsk = textureLoad(imageMsk, id.xy); - let colorDst = textureLoad(imageDst, id.xy); - var data: FragData; - data.Sc = colorSrc.rgb; - data.Sa = colorSrc.a; - data.Mc = colorMsk.rgb; - data.Ma = colorMsk.a; - data.Dc = colorDst.rgb; - data.Da = colorDst.a; - return data; -}; - -@compute @workgroup_size(8, 8) -fn cs_main_None(@builtin(global_invocation_id) id: vec3u) { - let d: FragData = getFragData(id.xy); - let Rc = d.Dc; - let Ra = d.Da; - textureStore(imageTgt, id.xy, vec4f(Rc, Ra)); -}; - -@compute @workgroup_size(8, 8) -fn cs_main_ClipPath(@builtin(global_invocation_id) id: vec3u) { - let d: FragData = getFragData(id.xy); - let Rc = d.Sc * d.Ma + d.Dc * (1.0 - d.Sa * d.Ma); - let Ra = d.Sa * d.Ma + d.Da * (1.0 - d.Sa * d.Ma); - textureStore(imageTgt, id.xy, vec4f(Rc, Ra)); -}; - -@compute @workgroup_size(8, 8) -fn cs_main_AlphaMask(@builtin(global_invocation_id) id: vec3u) { - let d: FragData = getFragData(id.xy); - let Rc = d.Sc * d.Ma + d.Dc * (1.0 - d.Sa * d.Ma); - let Ra = d.Sa * d.Ma + d.Da * (1.0 - d.Sa * d.Ma); - textureStore(imageTgt, id.xy, vec4f(Rc, Ra)); -}; - -@compute @workgroup_size(8, 8) -fn cs_main_InvAlphaMask(@builtin(global_invocation_id) id: vec3u) { - let d: FragData = getFragData(id.xy); - let Rc = d.Sc * (1.0 - d.Ma) + d.Dc * (1.0 - d.Sa * (1.0 - d.Ma)); - let Ra = d.Sa * (1.0 - d.Ma) + d.Da * (1.0 - d.Sa * (1.0 - d.Ma)); - textureStore(imageTgt, id.xy, vec4f(Rc, Ra)); -}; - -@compute @workgroup_size(8, 8) -fn cs_main_LumaMask(@builtin(global_invocation_id) id: vec3u) { - let d: FragData = getFragData(id.xy); - let luma: f32 = dot(d.Mc, vec3f(0.2125, 0.7154, 0.0721)); - let Rc = d.Sc * luma + d.Dc * (1.0 - d.Sa * luma); - let Ra = d.Sa * luma + d.Da * (1.0 - d.Sa * luma); - textureStore(imageTgt, id.xy, vec4f(Rc, Ra)); -}; - -@compute @workgroup_size(8, 8) -fn cs_main_InvLumaMask(@builtin(global_invocation_id) id: vec3u) { - let d: FragData = getFragData(id.xy); - let luma: f32 = dot(d.Mc, vec3f(0.2125, 0.7154, 0.0721)); - let Rc = d.Sc * (1.0 - luma) + d.Dc * (1.0 - d.Sa * (1.0 - luma)); - let Ra = d.Sa * (1.0 - luma) + d.Da * (1.0 - d.Sa * (1.0 - luma)); - textureStore(imageTgt, id.xy, vec4f(Rc, Ra)); -}; - -@compute @workgroup_size(8, 8) -fn cs_main_AddMask(@builtin(global_invocation_id) id: vec3u) { - let d: FragData = getFragData(id.xy); - let Rc = d.Sc; - let Ra = d.Sa + d.Ma * (1.0 - d.Sa); - textureStore(imageTgt, id.xy, vec4f(Rc, Ra)); -}; - -@compute @workgroup_size(8, 8) -fn cs_main_SubtractMask(@builtin(global_invocation_id) id: vec3u) { - let d: FragData = getFragData(id.xy); - let Rc = d.Sc; - let Ra = d.Sa * (1.0 - d.Ma); - textureStore(imageTgt, id.xy, vec4f(Rc, Ra)); -}; - -@compute @workgroup_size(8, 8) -fn cs_main_IntersectMask(@builtin(global_invocation_id) id: vec3u) { - let d: FragData = getFragData(id.xy); - let Rc = d.Sc; - let Ra = d.Sa * d.Ma; - textureStore(imageTgt, id.xy, vec4f(Rc, Ra)); -}; - -@compute @workgroup_size(8, 8) -fn cs_main_DifferenceMask(@builtin(global_invocation_id) id: vec3u) { - let d: FragData = getFragData(id.xy); - let Rc = d.Sc; - let Ra = d.Sa * (1.0 - d.Ma) + d.Ma * (1.0 - d.Sa); - textureStore(imageTgt, id.xy, vec4f(Rc, Ra)); -}; - -@compute @workgroup_size(8, 8) -fn cs_main_LightenMask(@builtin(global_invocation_id) id: vec3u) { - let d: FragData = getFragData(id.xy); - let Rc = d.Sc; - let Ra = max(d.Sa, d.Ma); - textureStore(imageTgt, id.xy, vec4f(Rc, Ra)); -}; - -@compute @workgroup_size(8, 8) -fn cs_main_DarkenMask(@builtin(global_invocation_id) id: vec3u) { - let d: FragData = getFragData(id.xy); - let Rc = d.Sc; - let Ra = min(d.Sa, d.Ma); - textureStore(imageTgt, id.xy, vec4f(Rc, Ra)); -}; -)"; - //************************************************************************ // compute shader source: copy //************************************************************************ diff --git a/src/renderer/wg_engine/tvgWgShaderSrc.h b/src/renderer/wg_engine/tvgWgShaderSrc.h index 85a1980d..2adfadb1 100755 --- a/src/renderer/wg_engine/tvgWgShaderSrc.h +++ b/src/renderer/wg_engine/tvgWgShaderSrc.h @@ -29,6 +29,7 @@ extern const char* cShaderSrc_Solid; extern const char* cShaderSrc_Linear; extern const char* cShaderSrc_Radial; extern const char* cShaderSrc_Image; +extern const char* cShaderSrc_Scene_Comp; // compute shader sources: blend, compose and merge path extern const char* cShaderSrc_MergeMasks; @@ -36,7 +37,6 @@ extern const char* cShaderSrc_BlendHeader_Solid; extern const char* cShaderSrc_BlendHeader_Gradient; extern const char* cShaderSrc_BlendHeader_Image; extern const char* cShaderSrc_Blend_Funcs; -extern const char* cShaderSrc_Compose; extern const char* cShaderSrc_Copy; #endif // _TVG_WG_SHEDER_SRC_H_