From 732a2be7e8d405ef744bd800a7dc181da977f18f Mon Sep 17 00:00:00 2001 From: Sergii Liebodkin Date: Fri, 23 Aug 2024 19:25:03 +0000 Subject: [PATCH] wg_engine: scene blending optimization - used hardware blending stage for scene blending - used AABB for scene blending - reduced number of offfscreen buffers coping - reduced number of render pass switching - used render pipelines abilities to convert offscreen pixel format to screen pixel format - removed unused shaders --- src/renderer/wg_engine/tvgWgBindGroups.cpp | 20 ------ src/renderer/wg_engine/tvgWgBindGroups.h | 2 - src/renderer/wg_engine/tvgWgCommon.cpp | 2 +- src/renderer/wg_engine/tvgWgCompositor.cpp | 51 +++++++++++++--- src/renderer/wg_engine/tvgWgCompositor.h | 3 + src/renderer/wg_engine/tvgWgPipelines.cpp | 71 +++++++++++++++------- src/renderer/wg_engine/tvgWgPipelines.h | 11 ++-- src/renderer/wg_engine/tvgWgRenderer.cpp | 46 ++++++-------- src/renderer/wg_engine/tvgWgRenderer.h | 5 -- src/renderer/wg_engine/tvgWgShaderSrc.cpp | 66 +++++++++++++++----- src/renderer/wg_engine/tvgWgShaderSrc.h | 3 +- 11 files changed, 174 insertions(+), 106 deletions(-) diff --git a/src/renderer/wg_engine/tvgWgBindGroups.cpp b/src/renderer/wg_engine/tvgWgBindGroups.cpp index 4478d8d2..58f8061c 100755 --- a/src/renderer/wg_engine/tvgWgBindGroups.cpp +++ b/src/renderer/wg_engine/tvgWgBindGroups.cpp @@ -45,15 +45,6 @@ WGPUBindGroup WgBindGroupLayouts::createBindGroupTexSampledBuff1Un(WGPUSampler s } -WGPUBindGroup WgBindGroupLayouts::createBindGroupScreen1WO(WGPUTextureView texView) { - const WGPUBindGroupEntry bindGroupEntrys[] = { - { .binding = 0, .textureView = texView } - }; - const WGPUBindGroupDescriptor bindGroupDesc { .layout = layoutTexScreen1WO, .entryCount = 1, .entries = bindGroupEntrys }; - return wgpuDeviceCreateBindGroup(device, &bindGroupDesc); -} - - WGPUBindGroup WgBindGroupLayouts::createBindGroupStrorage1WO(WGPUTextureView texView) { const WGPUBindGroupEntry bindGroupEntrys[] = { @@ -150,7 +141,6 @@ void WgBindGroupLayouts::initialize(WgContext& context) const WGPUTextureBindingLayout texture = { .sampleType = WGPUTextureSampleType_Float, .viewDimension = WGPUTextureViewDimension_2D }; const WGPUStorageTextureBindingLayout storageTextureWO { .access = WGPUStorageTextureAccess_WriteOnly, .format = WGPUTextureFormat_RGBA8Unorm, .viewDimension = WGPUTextureViewDimension_2D }; const WGPUStorageTextureBindingLayout storageTextureRO { .access = WGPUStorageTextureAccess_ReadOnly, .format = WGPUTextureFormat_RGBA8Unorm, .viewDimension = WGPUTextureViewDimension_2D }; - const WGPUStorageTextureBindingLayout storageScreenWO { .access = WGPUStorageTextureAccess_WriteOnly, .format = WGPUTextureFormat_BGRA8Unorm, .viewDimension = WGPUTextureViewDimension_2D }; const WGPUBufferBindingLayout bufferUniform { .type = WGPUBufferBindingType_Uniform }; { // bind group layout tex sampled @@ -174,15 +164,6 @@ void WgBindGroupLayouts::initialize(WgContext& context) assert(layoutTexSampledBuff1Un); } - { // bind group layout tex screen 1 RO - const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { - { .binding = 0, .visibility = visibility_frag, .storageTexture = storageScreenWO } - }; - const WGPUBindGroupLayoutDescriptor bindGroupLayoutDesc { .entryCount = 1, .entries = bindGroupLayoutEntries }; - layoutTexScreen1WO = wgpuDeviceCreateBindGroupLayout(device, &bindGroupLayoutDesc); - assert(layoutTexScreen1WO); - } - { // bind group layout tex storage 1 WO const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { { .binding = 0, .visibility = visibility_frag, .storageTexture = storageTextureWO } @@ -259,7 +240,6 @@ void WgBindGroupLayouts::release(WgContext& context) wgpuBindGroupLayoutRelease(layoutBuffer3Un); wgpuBindGroupLayoutRelease(layoutBuffer2Un); wgpuBindGroupLayoutRelease(layoutBuffer1Un); - wgpuBindGroupLayoutRelease(layoutTexScreen1WO); wgpuBindGroupLayoutRelease(layoutTexStrorage3RO); wgpuBindGroupLayoutRelease(layoutTexStrorage2RO); wgpuBindGroupLayoutRelease(layoutTexStrorage1RO); diff --git a/src/renderer/wg_engine/tvgWgBindGroups.h b/src/renderer/wg_engine/tvgWgBindGroups.h index 88588dec..82a31eff 100755 --- a/src/renderer/wg_engine/tvgWgBindGroups.h +++ b/src/renderer/wg_engine/tvgWgBindGroups.h @@ -31,7 +31,6 @@ private: public: WGPUBindGroupLayout layoutTexSampled{}; WGPUBindGroupLayout layoutTexSampledBuff1Un{}; - WGPUBindGroupLayout layoutTexScreen1WO{}; WGPUBindGroupLayout layoutTexStrorage1WO{}; WGPUBindGroupLayout layoutTexStrorage1RO{}; WGPUBindGroupLayout layoutTexStrorage2RO{}; @@ -42,7 +41,6 @@ public: public: WGPUBindGroup createBindGroupTexSampled(WGPUSampler sampler, WGPUTextureView texView); WGPUBindGroup createBindGroupTexSampledBuff1Un(WGPUSampler sampler, WGPUTextureView texView, WGPUBuffer buff); - WGPUBindGroup createBindGroupScreen1WO(WGPUTextureView texView); WGPUBindGroup createBindGroupStrorage1WO(WGPUTextureView texView); WGPUBindGroup createBindGroupStrorage1RO(WGPUTextureView texView); WGPUBindGroup createBindGroupStrorage2RO(WGPUTextureView texView0, WGPUTextureView texView1); diff --git a/src/renderer/wg_engine/tvgWgCommon.cpp b/src/renderer/wg_engine/tvgWgCommon.cpp index c97cf703..3e5af9a3 100755 --- a/src/renderer/wg_engine/tvgWgCommon.cpp +++ b/src/renderer/wg_engine/tvgWgCommon.cpp @@ -48,7 +48,7 @@ void WgContext::initialize(WGPUInstance instance, WGPUSurface surface) // get adapter and surface properties WGPUFeatureName featureNames[32]{}; size_t featuresCount = wgpuAdapterEnumerateFeatures(adapter, featureNames); - preferredFormat = wgpuSurfaceGetPreferredFormat(surface, adapter); + preferredFormat = WGPUTextureFormat_BGRA8Unorm; // request device const WGPUDeviceDescriptor deviceDesc { .nextInChain = nullptr, .label = "The device", .requiredFeatureCount = featuresCount, .requiredFeatures = featureNames }; diff --git a/src/renderer/wg_engine/tvgWgCompositor.cpp b/src/renderer/wg_engine/tvgWgCompositor.cpp index 3b17467b..d30c188e 100755 --- a/src/renderer/wg_engine/tvgWgCompositor.cpp +++ b/src/renderer/wg_engine/tvgWgCompositor.cpp @@ -87,6 +87,17 @@ WgPipelineBlendType WgCompositor::blendMethodToBlendType(BlendMethod blendMethod } +RenderRegion WgCompositor::shrinkRenderRegion(RenderRegion& rect) +{ + // cut viewport to screen dimensions + int32_t xmin = std::max(0, std::min((int32_t)width, rect.x)); + int32_t ymin = std::max(0, std::min((int32_t)height, rect.y)); + int32_t xmax = std::max(0, std::min((int32_t)width, rect.x + rect.w)); + int32_t ymax = std::max(0, std::min((int32_t)height, rect.y + rect.h)); + return { xmin, ymin, xmax - xmin, ymax - ymin }; +} + + void WgCompositor::beginRenderPass(WGPUCommandEncoder commandEncoder, WgRenderStorage* target, bool clear) { assert(commandEncoder); @@ -258,7 +269,19 @@ void WgCompositor::blendImage(WgContext& context, WgRenderDataPicture* renderDat }; -// TODO: use direct mask applience +void WgCompositor::blendScene(WgContext& context, WgRenderStorage* src, WgCompose* cmp) +{ + assert(currentTarget); + RenderRegion rect = shrinkRenderRegion(cmp->aabb); + wgpuRenderPassEncoderSetScissorRect(renderPassEncoder, rect.x, rect.y, rect.w, rect.h); + wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 0); + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 0, src->bindGroupTexure, 0, nullptr); + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 1, bindGroupOpacities[cmp->opacity], 0, nullptr); + wgpuRenderPassEncoderSetPipeline(renderPassEncoder, pipelines->sceneBlend); + meshData.drawImage(context, renderPassEncoder); +} + + void WgCompositor::composeShape(WgContext& context, WgRenderDataShape* renderData, WgRenderStorage* mask, CompositeMethod composeMethod) { assert(mask); @@ -282,7 +305,6 @@ void WgCompositor::composeShape(WgContext& context, WgRenderDataShape* renderDat } -// TODO: use direct mask applience void WgCompositor::composeStrokes(WgContext& context, WgRenderDataShape* renderData, WgRenderStorage* mask, CompositeMethod composeMethod) { assert(mask); @@ -306,7 +328,6 @@ void WgCompositor::composeStrokes(WgContext& context, WgRenderDataShape* renderD } -// TODO: use direct mask applience void WgCompositor::composeImage(WgContext& context, WgRenderDataPicture* renderData, WgRenderStorage* mask, CompositeMethod composeMethod) { assert(mask); @@ -333,13 +354,7 @@ void WgCompositor::composeScene(WgContext& context, WgRenderStorage* src, WgRend 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 }; + RenderRegion rect = shrinkRenderRegion(cmp->aabb); composeRegion(context, src, mask, cmp->method, rect); } @@ -520,3 +535,19 @@ void WgCompositor::blend(WGPUCommandEncoder encoder, WgRenderStorage* src, WgRen wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, (width + 7) / 8, (height + 7) / 8, 1); wgpuComputePassEncoderEnd(computePassEncoder); } + + +void WgCompositor::blit(WgContext& context, WGPUCommandEncoder encoder, WgRenderStorage* src, WGPUTextureView dstView) { + WGPURenderPassDepthStencilAttachment depthStencilAttachment{ .view = texViewStencil, .stencilLoadOp = WGPULoadOp_Load, .stencilStoreOp = WGPUStoreOp_Discard }; + WGPURenderPassColorAttachment colorAttachment { .view = dstView, .loadOp = WGPULoadOp_Load, .storeOp = WGPUStoreOp_Store }; + #ifdef __EMSCRIPTEN__ + colorAttachment.depthSlice = WGPU_DEPTH_SLICE_UNDEFINED; + #endif + WGPURenderPassDescriptor renderPassDesc{ .colorAttachmentCount = 1, .colorAttachments = &colorAttachment, .depthStencilAttachment = &depthStencilAttachment }; + WGPURenderPassEncoder renderPass = wgpuCommandEncoderBeginRenderPass(encoder, &renderPassDesc); + wgpuRenderPassEncoderSetBindGroup(renderPass, 0, src->bindGroupTexure, 0, nullptr); + wgpuRenderPassEncoderSetPipeline(renderPass, pipelines->blit); + meshData.drawImage(context, renderPass); + wgpuRenderPassEncoderEnd(renderPass); + wgpuRenderPassEncoderRelease(renderPass); +} diff --git a/src/renderer/wg_engine/tvgWgCompositor.h b/src/renderer/wg_engine/tvgWgCompositor.h index 5905c36c..574d8ed1 100755 --- a/src/renderer/wg_engine/tvgWgCompositor.h +++ b/src/renderer/wg_engine/tvgWgCompositor.h @@ -58,6 +58,7 @@ private: static WgPipelineBlendType blendMethodToBlendType(BlendMethod blendMethod); void composeRegion(WgContext& context, WgRenderStorage* src, WgRenderStorage* mask, CompositeMethod composeMethod, RenderRegion& rect); + RenderRegion shrinkRenderRegion(RenderRegion& rect); public: // render target dimensions uint32_t width{}; @@ -77,6 +78,7 @@ public: void blendShape(WgContext& context, WgRenderDataShape* renderData, BlendMethod blendMethod); void blendStrokes(WgContext& context, WgRenderDataShape* renderData, BlendMethod blendMethod); void blendImage(WgContext& context, WgRenderDataPicture* renderData, BlendMethod blendMethod); + void blendScene(WgContext& context, WgRenderStorage* src, WgCompose* cmp); void composeShape(WgContext& context, WgRenderDataShape* renderData, WgRenderStorage* mask, CompositeMethod composeMethod); void composeStrokes(WgContext& context, WgRenderDataShape* renderData, WgRenderStorage* mask, CompositeMethod composeMethod); @@ -90,6 +92,7 @@ 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 blit(WgContext& context, WGPUCommandEncoder encoder, WgRenderStorage* src, WGPUTextureView dstView); }; #endif // _TVG_WG_COMPOSITOR_H_ diff --git a/src/renderer/wg_engine/tvgWgPipelines.cpp b/src/renderer/wg_engine/tvgWgPipelines.cpp index 260a0889..504cfab1 100755 --- a/src/renderer/wg_engine/tvgWgPipelines.cpp +++ b/src/renderer/wg_engine/tvgWgPipelines.cpp @@ -46,12 +46,12 @@ WGPURenderPipeline WgPipelines::createRenderPipeline( const WGPUShaderModule shaderModule, const char* vsEntryPoint, const char* fsEntryPoint, const WGPUPipelineLayout pipelineLayout, const WGPUVertexBufferLayout *vertexBufferLayouts, const uint32_t vertexBufferLayoutsCount, - const WGPUColorWriteMaskFlags writeMask, + const WGPUColorWriteMaskFlags writeMask, const WGPUTextureFormat colorTargetFormat, const WGPUCompareFunction stencilFunctionFrnt, const WGPUStencilOperation stencilOperationFrnt, const WGPUCompareFunction stencilFunctionBack, const WGPUStencilOperation stencilOperationBack, const WGPUPrimitiveState primitiveState, const WGPUMultisampleState multisampleState, const WGPUBlendState blendState) { - const WGPUColorTargetState colorTargetState { .format = WGPUTextureFormat_RGBA8Unorm, .blend = &blendState, .writeMask = writeMask }; + const WGPUColorTargetState colorTargetState { .format = colorTargetFormat, .blend = &blendState, .writeMask = writeMask }; const WGPUColorTargetState colorTargetStates[] { colorTargetState }; const WGPUDepthStencilState depthStencilState { .format = WGPUTextureFormat_Stencil8, .depthCompare = WGPUCompareFunction_Always, @@ -147,6 +147,7 @@ void WgPipelines::initialize(WgContext& context) const WGPUVertexBufferLayout vertexBufferLayoutsImage[] { vertexBufferLayoutPos, vertexBufferLayoutTex }; const WGPUPrimitiveState primitiveState { .topology = WGPUPrimitiveTopology_TriangleList }; const WGPUMultisampleState multisampleState { .count = 1, .mask = 0xFFFFFFFF, .alphaToCoverageEnabled = false }; + const WGPUTextureFormat offscreenTargetFormat = WGPUTextureFormat_RGBA8Unorm; // blend states const WGPUBlendState blendStateSrc { @@ -187,6 +188,13 @@ void WgPipelines::initialize(WgContext& context) layouts.layoutTexSampled, layouts.layoutTexSampled }; + const WGPUBindGroupLayout bindGroupLayoutsSceneBlend[] = { + layouts.layoutTexSampled, + layouts.layoutBuffer1Un + }; + const WGPUBindGroupLayout bindGroupLayoutsBlit[] = { + layouts.layoutTexSampled + }; const WGPUBindGroupLayout bindGroupLayoutsMergeMasks[] = { layouts.layoutTexStrorage1RO, layouts.layoutTexStrorage1RO, @@ -198,10 +206,6 @@ void WgPipelines::initialize(WgContext& context) layouts.layoutTexStrorage1WO, layouts.layoutBuffer1Un }; - const WGPUBindGroupLayout bindGroupLayoutsCopy[] = { - layouts.layoutTexStrorage1RO, - layouts.layoutTexScreen1WO - }; // pipeline layouts layoutStencil = createPipelineLayout(context.device, bindGroupLayoutsStencil, 2); @@ -209,9 +213,10 @@ void WgPipelines::initialize(WgContext& context) layoutGradient = createPipelineLayout(context.device, bindGroupLayoutsGradient, 3); layoutImage = createPipelineLayout(context.device, bindGroupLayoutsImage, 3); layoutSceneComp = createPipelineLayout(context.device, bindGroupLayoutsSceneComp, 2); + layoutSceneBlend = createPipelineLayout(context.device, bindGroupLayoutsSceneBlend, 2); + layoutBlit = createPipelineLayout(context.device, bindGroupLayoutsBlit, 1); layoutBlend = createPipelineLayout(context.device, bindGroupLayoutsBlend, 4); layoutMergeMasks = createPipelineLayout(context.device, bindGroupLayoutsMergeMasks, 3); - layoutCopy = createPipelineLayout(context.device, bindGroupLayoutsCopy, 2); // graphics shader modules shaderStencil = createShaderModule(context.device, "The shader stencil", cShaderSrc_Stencil); @@ -220,16 +225,17 @@ void WgPipelines::initialize(WgContext& context) 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); + shaderSceneBlend = createShaderModule(context.device, "The shader scene blend", cShaderSrc_Scene_Blend); + shaderBlit = createShaderModule(context.device, "The shader blit", cShaderSrc_Blit); // computes shader modules shaderMergeMasks = createShaderModule(context.device, "The shader merge mask", cShaderSrc_MergeMasks); - shaderCopy = createShaderModule(context.device, "The shader copy", cShaderSrc_Copy); // render pipeline winding winding = createRenderPipeline( context.device, "The render pipeline winding", shaderStencil, "vs_main", "fs_main", layoutStencil, vertexBufferLayoutsShape, 1, - WGPUColorWriteMask_None, + WGPUColorWriteMask_None, offscreenTargetFormat, WGPUCompareFunction_Always, WGPUStencilOperation_IncrementWrap, WGPUCompareFunction_Always, WGPUStencilOperation_DecrementWrap, primitiveState, multisampleState, blendStateSrc); @@ -238,7 +244,7 @@ void WgPipelines::initialize(WgContext& context) context.device, "The render pipeline even-odd", shaderStencil, "vs_main", "fs_main", layoutStencil, vertexBufferLayoutsShape, 1, - WGPUColorWriteMask_None, + WGPUColorWriteMask_None, offscreenTargetFormat, WGPUCompareFunction_Always, WGPUStencilOperation_Invert, WGPUCompareFunction_Always, WGPUStencilOperation_Invert, primitiveState, multisampleState, blendStateSrc); @@ -247,7 +253,7 @@ void WgPipelines::initialize(WgContext& context) context.device, "The render pipeline direct", shaderStencil,"vs_main", "fs_main", layoutStencil, vertexBufferLayoutsShape, 1, - WGPUColorWriteMask_None, + WGPUColorWriteMask_None, offscreenTargetFormat, WGPUCompareFunction_Always, WGPUStencilOperation_Replace, WGPUCompareFunction_Always, WGPUStencilOperation_Replace, primitiveState, multisampleState, blendStateSrc); @@ -256,7 +262,7 @@ void WgPipelines::initialize(WgContext& context) context.device, "The render pipeline clip path", shaderStencil, "vs_main", "fs_main", layoutStencil, vertexBufferLayoutsShape, 1, - WGPUColorWriteMask_All, + WGPUColorWriteMask_All, offscreenTargetFormat, WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero, WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero, primitiveState, multisampleState, blendStateSrc); @@ -267,7 +273,7 @@ void WgPipelines::initialize(WgContext& context) context.device, "The render pipeline solid", shaderSolid, "vs_main", "fs_main", layoutSolid, vertexBufferLayoutsShape, 1, - WGPUColorWriteMask_All, + WGPUColorWriteMask_All, offscreenTargetFormat, WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero, WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero, primitiveState, multisampleState, blendStates[i]); @@ -279,7 +285,7 @@ void WgPipelines::initialize(WgContext& context) context.device, "The render pipeline radial", shaderRadial, "vs_main", "fs_main", layoutGradient, vertexBufferLayoutsShape, 1, - WGPUColorWriteMask_All, + WGPUColorWriteMask_All, offscreenTargetFormat, WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero, WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero, primitiveState, multisampleState, blendStates[i]); @@ -291,7 +297,7 @@ void WgPipelines::initialize(WgContext& context) context.device, "The render pipeline linear", shaderLinear, "vs_main", "fs_main", layoutGradient, vertexBufferLayoutsShape, 1, - WGPUColorWriteMask_All, + WGPUColorWriteMask_All, offscreenTargetFormat, WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero, WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero, primitiveState, multisampleState, blendStates[i]); @@ -303,13 +309,32 @@ void WgPipelines::initialize(WgContext& context) context.device, "The render pipeline image", shaderImage, "vs_main", "fs_main", layoutImage, vertexBufferLayoutsImage, 2, - WGPUColorWriteMask_All, + WGPUColorWriteMask_All, offscreenTargetFormat, WGPUCompareFunction_Always, WGPUStencilOperation_Zero, WGPUCompareFunction_Always, WGPUStencilOperation_Zero, primitiveState, multisampleState, blendStates[i]); } - // compute shader names + // render pipeline blit + blit = createRenderPipeline(context.device, "The render pipeline blit", + shaderBlit, "vs_main", "fs_main", layoutBlit, + vertexBufferLayoutsImage, 2, + // must be preferred screen pixel format + WGPUColorWriteMask_All, context.preferredFormat, + WGPUCompareFunction_Always, WGPUStencilOperation_Zero, + WGPUCompareFunction_Always, WGPUStencilOperation_Zero, + primitiveState, multisampleState, blendStateSrc); + + // render pipeline blit + sceneBlend = createRenderPipeline(context.device, "The render pipeline scene blend", + shaderSceneBlend, "vs_main", "fs_main", layoutSceneBlend, + vertexBufferLayoutsImage, 2, + WGPUColorWriteMask_All, offscreenTargetFormat, + WGPUCompareFunction_Always, WGPUStencilOperation_Zero, + WGPUCompareFunction_Always, WGPUStencilOperation_Zero, + primitiveState, multisampleState, blendStateNrm); + + // compose shader names const char* shaderComposeNames[] { "fs_main_None", "fs_main_ClipPath", @@ -352,7 +377,7 @@ void WgPipelines::initialize(WgContext& context) context.device, "The render pipeline scene composition", shaderSceneComp, "vs_main", shaderComposeNames[i], layoutSceneComp, vertexBufferLayoutsImage, 2, - WGPUColorWriteMask_All, + WGPUColorWriteMask_All, offscreenTargetFormat, WGPUCompareFunction_Always, WGPUStencilOperation_Zero, WGPUCompareFunction_Always, WGPUStencilOperation_Zero, primitiveState, multisampleState, composeBlends[i]); @@ -360,7 +385,6 @@ void WgPipelines::initialize(WgContext& context) // 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); // compute shader blend names const char* shaderBlendNames[] { @@ -399,7 +423,9 @@ void WgPipelines::initialize(WgContext& context) void WgPipelines::releaseGraphicHandles(WgContext& context) { + releaseRenderPipeline(blit); releaseRenderPipeline(clipPath); + releaseRenderPipeline(sceneBlend); size_t pipesSceneCompCnt = sizeof(sceneComp) / sizeof(sceneComp[0]); for (uint32_t i = 0; i < pipesSceneCompCnt; i++) releaseRenderPipeline(sceneComp[i]); @@ -412,11 +438,15 @@ void WgPipelines::releaseGraphicHandles(WgContext& context) releaseRenderPipeline(direct); releaseRenderPipeline(evenodd); releaseRenderPipeline(winding); + releasePipelineLayout(layoutBlit); + releasePipelineLayout(layoutSceneBlend); releasePipelineLayout(layoutSceneComp); releasePipelineLayout(layoutImage); releasePipelineLayout(layoutGradient); releasePipelineLayout(layoutSolid); releasePipelineLayout(layoutStencil); + releaseShaderModule(shaderBlit); + releaseShaderModule(shaderSceneBlend); releaseShaderModule(shaderSceneComp); releaseShaderModule(shaderImage); releaseShaderModule(shaderLinear); @@ -434,16 +464,13 @@ void WgPipelines::releaseComputeHandles(WgContext& context) releaseComputePipeline(blendSolid[i]); releaseComputePipeline(blendGradient[i]); } - releaseComputePipeline(copy); releaseComputePipeline(mergeMasks); releasePipelineLayout(layoutBlend); releasePipelineLayout(layoutMergeMasks); - releasePipelineLayout(layoutCopy); releaseShaderModule(shaderBlendImage); releaseShaderModule(shaderBlendGradient); releaseShaderModule(shaderBlendSolid); releaseShaderModule(shaderMergeMasks); - releaseShaderModule(shaderCopy); } void WgPipelines::release(WgContext& context) diff --git a/src/renderer/wg_engine/tvgWgPipelines.h b/src/renderer/wg_engine/tvgWgPipelines.h index 57320b45..8a02a2f6 100755 --- a/src/renderer/wg_engine/tvgWgPipelines.h +++ b/src/renderer/wg_engine/tvgWgPipelines.h @@ -36,12 +36,13 @@ private: WGPUShaderModule shaderLinear{}; WGPUShaderModule shaderImage{}; WGPUShaderModule shaderSceneComp{}; + WGPUShaderModule shaderSceneBlend{}; + WGPUShaderModule shaderBlit{}; // compute pipeline shaders WGPUShaderModule shaderMergeMasks; WGPUShaderModule shaderBlendSolid; WGPUShaderModule shaderBlendGradient; WGPUShaderModule shaderBlendImage; - WGPUShaderModule shaderCopy; private: // graphics pipeline layouts WGPUPipelineLayout layoutStencil{}; @@ -49,10 +50,11 @@ private: WGPUPipelineLayout layoutGradient{}; WGPUPipelineLayout layoutImage{}; WGPUPipelineLayout layoutSceneComp{}; + WGPUPipelineLayout layoutSceneBlend{}; + WGPUPipelineLayout layoutBlit{}; // compute pipeline layouts WGPUPipelineLayout layoutMergeMasks{}; WGPUPipelineLayout layoutBlend{}; - WGPUPipelineLayout layoutCopy{}; public: // graphics pipeline WgBindGroupLayouts layouts; @@ -64,13 +66,14 @@ public: WGPURenderPipeline linear[3]{}; WGPURenderPipeline image[3]{}; WGPURenderPipeline sceneComp[12]; + WGPURenderPipeline sceneBlend; + WGPURenderPipeline blit{}; WGPURenderPipeline clipPath{}; // compute pipeline WGPUComputePipeline mergeMasks; WGPUComputePipeline blendSolid[14]; WGPUComputePipeline blendGradient[14]; WGPUComputePipeline blendImage[14]; - WGPUComputePipeline copy; private: void releaseGraphicHandles(WgContext& context); void releaseComputeHandles(WgContext& context); @@ -82,7 +85,7 @@ private: const WGPUShaderModule shaderModule, const char* vsEntryPoint, const char* fsEntryPoint, const WGPUPipelineLayout pipelineLayout, const WGPUVertexBufferLayout *vertexBufferLayouts, const uint32_t vertexBufferLayoutsCount, - const WGPUColorWriteMaskFlags writeMask, + const WGPUColorWriteMaskFlags writeMask, const WGPUTextureFormat colorTargetFormat, const WGPUCompareFunction stencilFunctionFrnt, const WGPUStencilOperation stencilOperationFrnt, const WGPUCompareFunction stencilFunctionBack, const WGPUStencilOperation stencilOperationBack, const WGPUPrimitiveState primitiveState, const WGPUMultisampleState multisampleState, const WGPUBlendState blendState); diff --git a/src/renderer/wg_engine/tvgWgRenderer.cpp b/src/renderer/wg_engine/tvgWgRenderer.cpp index 7c0fab52..35d3a88a 100755 --- a/src/renderer/wg_engine/tvgWgRenderer.cpp +++ b/src/renderer/wg_engine/tvgWgRenderer.cpp @@ -48,9 +48,6 @@ void WgRenderer::initialize() void WgRenderer::release() { disposeObjects(); - mContext.releaseTexture(mTexScreen); - mContext.releaseTextureView(mTexViewScreen); - mContext.pipelines->layouts.releaseBindGroup(mBindGroupScreen); mStorageRoot.release(mContext); mRenderStoragePool.release(mContext); mRenderDataShapePool.release(mContext); @@ -252,25 +249,14 @@ bool WgRenderer::sync() // get current texture WGPUSurfaceTexture surfaceTexture{}; wgpuSurfaceGetCurrentTexture(mContext.surface, &surfaceTexture); + WGPUTextureView dstView = mContext.createTextureView(surfaceTexture.texture); // create command encoder const WGPUCommandEncoderDescriptor commandEncoderDesc{}; WGPUCommandEncoder commandEncoder = wgpuDeviceCreateCommandEncoder(mContext.device, &commandEncoderDesc); - // copy render to screen with conversion rgba to bgra (screen format) - const WGPUComputePassDescriptor computePassDescriptor{}; - WGPUComputePassEncoder computePassEncoder = wgpuCommandEncoderBeginComputePass(commandEncoder, &computePassDescriptor); - wgpuComputePassEncoderSetBindGroup(computePassEncoder, 0, mStorageRoot.bindGroupRead, 0, nullptr); - wgpuComputePassEncoderSetBindGroup(computePassEncoder, 1, mBindGroupScreen, 0, nullptr); - wgpuComputePassEncoderSetPipeline(computePassEncoder, mContext.pipelines->copy); - wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, (mStorageRoot.width + 7) / 8, (mStorageRoot.height + 7) / 8, 1); - wgpuComputePassEncoderEnd(computePassEncoder); - - // copy dst storage to temporary read only storage - const WGPUImageCopyTexture texSrc { .texture = mTexScreen }; - const WGPUImageCopyTexture texDst { .texture = surfaceTexture.texture }; - const WGPUExtent3D copySize { .width = mTargetSurface.w, .height = mTargetSurface.h, .depthOrArrayLayers = 1 }; - wgpuCommandEncoderCopyTextureToTexture(commandEncoder, &texSrc, &texDst, ©Size); + // show root offscreen buffer + mCompositor.blit(mContext, commandEncoder, &mStorageRoot, dstView); // release command encoder const WGPUCommandBufferDescriptor commandBufferDesc{}; @@ -278,6 +264,7 @@ bool WgRenderer::sync() wgpuQueueSubmit(mContext.queue, 1, &commandsBuffer); wgpuCommandBufferRelease(commandsBuffer); wgpuCommandEncoderRelease(commandEncoder); + mContext.releaseTextureView(dstView); return true; } @@ -295,8 +282,8 @@ bool WgRenderer::target(WGPUInstance instance, WGPUSurface surface, uint32_t w, WGPUSurfaceConfiguration surfaceConfiguration { .device = mContext.device, - .format = WGPUTextureFormat_BGRA8Unorm, - .usage = WGPUTextureUsage_CopyDst, + .format = mContext.preferredFormat, + .usage = WGPUTextureUsage_RenderAttachment, .width = w, .height = h, #ifdef __EMSCRIPTEN__ .presentMode = WGPUPresentMode_Fifo, @@ -309,10 +296,6 @@ bool WgRenderer::target(WGPUInstance instance, WGPUSurface surface, uint32_t w, mRenderStoragePool.initialize(mContext, w, h); mStorageRoot.initialize(mContext, w, h); mCompositor.initialize(mContext, w, h); - // screen buffer - mTexScreen = mContext.createTexStorage(w, h, WGPUTextureFormat_BGRA8Unorm); - mTexViewScreen = mContext.createTextureView(mTexScreen); - mBindGroupScreen = mContext.pipelines->layouts.createBindGroupScreen1WO(mTexViewScreen); return true; } @@ -355,12 +338,21 @@ bool WgRenderer::endComposite(Compositor* cmp) WgRenderStorage* src = mRenderStorageStack.last(); mRenderStorageStack.pop(); WgRenderStorage* dst = mRenderStorageStack.last(); - // apply blend - mCompositor.blend(mCommandEncoder, src, dst, comp->opacity, comp->blend, WgRenderRasterType::Image); + // apply normal blend + if (comp->blend == BlendMethod::Normal) { + // begin previous render pass + mCompositor.beginRenderPass(mCommandEncoder, dst, false); + // apply blend + mCompositor.blendScene(mContext, src, comp); + // apply custom blend + } else { + // apply custom blend + mCompositor.blend(mCommandEncoder, src, dst, comp->opacity, comp->blend, WgRenderRasterType::Image); + // begin previous render pass + mCompositor.beginRenderPass(mCommandEncoder, dst, false); + } // 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(); diff --git a/src/renderer/wg_engine/tvgWgRenderer.h b/src/renderer/wg_engine/tvgWgRenderer.h index f077df55..d90d2a68 100755 --- a/src/renderer/wg_engine/tvgWgRenderer.h +++ b/src/renderer/wg_engine/tvgWgRenderer.h @@ -75,11 +75,6 @@ private: WgPipelines mPipelines; WgCompositor mCompositor; - // screen buffer - WGPUTexture mTexScreen{}; - WGPUTextureView mTexViewScreen{}; - WGPUBindGroup mBindGroupScreen{}; - Surface mTargetSurface; BlendMethod mBlendMethod{}; RenderRegion mViewport{}; diff --git a/src/renderer/wg_engine/tvgWgShaderSrc.cpp b/src/renderer/wg_engine/tvgWgShaderSrc.cpp index bb191720..a3114a6c 100755 --- a/src/renderer/wg_engine/tvgWgShaderSrc.cpp +++ b/src/renderer/wg_engine/tvgWgShaderSrc.cpp @@ -291,6 +291,58 @@ fn fs_main_DarkenMask(in: VertexOutput) -> @location(0) vec4f { }; )"; + +//************************************************************************ +// graphics shader source: scene blend +//************************************************************************ + +const char* cShaderSrc_Scene_Blend = 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 So : f32; + +@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(in: VertexOutput) -> @location(0) vec4f { + return textureSample(uTextureSrc, uSamplerSrc, in.texCoord.xy) * So; +}; +)"; + +//************************************************************************ +// graphics shader source: texture blit +//************************************************************************ + +const char* cShaderSrc_Blit = 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; + +@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(in: VertexOutput) -> @location(0) vec4f { + return textureSample(uTextureSrc, uSamplerSrc, in.texCoord.xy); +}; +)"; + //************************************************************************ // compute shader source: merge clip path masks //************************************************************************ @@ -518,17 +570,3 @@ fn cs_main_SoftLight(@builtin(global_invocation_id) id: vec3u) { textureStore(imageTgt, id.xy, postProcess(d, vec4f(Rc, 1.0))); }; )"; - -//************************************************************************ -// compute shader source: copy -//************************************************************************ - -const char* cShaderSrc_Copy = R"( -@group(0) @binding(0) var imageSrc : texture_storage_2d; -@group(1) @binding(0) var imageDst : texture_storage_2d; - -@compute @workgroup_size(8, 8) -fn cs_main(@builtin(global_invocation_id) id: vec3u) { - textureStore(imageDst, id.xy, textureLoad(imageSrc, id.xy)); -}; -)"; diff --git a/src/renderer/wg_engine/tvgWgShaderSrc.h b/src/renderer/wg_engine/tvgWgShaderSrc.h index 2adfadb1..6e2c488e 100755 --- a/src/renderer/wg_engine/tvgWgShaderSrc.h +++ b/src/renderer/wg_engine/tvgWgShaderSrc.h @@ -30,6 +30,8 @@ extern const char* cShaderSrc_Linear; extern const char* cShaderSrc_Radial; extern const char* cShaderSrc_Image; extern const char* cShaderSrc_Scene_Comp; +extern const char* cShaderSrc_Scene_Blend; +extern const char* cShaderSrc_Blit; // compute shader sources: blend, compose and merge path extern const char* cShaderSrc_MergeMasks; @@ -37,6 +39,5 @@ 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_Copy; #endif // _TVG_WG_SHEDER_SRC_H_