diff --git a/src/renderer/wg_engine/tvgWgBindGroups.cpp b/src/renderer/wg_engine/tvgWgBindGroups.cpp index 3e194b16..1341074f 100644 --- a/src/renderer/wg_engine/tvgWgBindGroups.cpp +++ b/src/renderer/wg_engine/tvgWgBindGroups.cpp @@ -32,10 +32,12 @@ WGPUBindGroupLayout WgBindGroupLinearGradient::layout = nullptr; WGPUBindGroupLayout WgBindGroupRadialGradient::layout = nullptr; WGPUBindGroupLayout WgBindGroupPicture::layout = nullptr; // composition and blending properties gropus -WGPUBindGroupLayout WgBindGroupOpacity::layout = nullptr; WGPUBindGroupLayout WgBindGroupTexture::layout = nullptr; -WGPUBindGroupLayout WgBindGroupStorageTexture::layout = nullptr; +WGPUBindGroupLayout WgBindGroupTextureStorage::layout = nullptr; WGPUBindGroupLayout WgBindGroupTextureSampled::layout = nullptr; +WGPUBindGroupLayout WgBindGroupOpacity::layout = nullptr; +WGPUBindGroupLayout WgBindGroupBlendMethod::layout = nullptr; +WGPUBindGroupLayout WgBindGroupCompositeMethod::layout = nullptr; WGPUBindGroupLayout WgBindGroupCanvas::getLayout(WGPUDevice device) @@ -264,50 +266,6 @@ void WgBindGroupPicture::release() } -WGPUBindGroupLayout WgBindGroupOpacity::getLayout(WGPUDevice device) -{ - if (layout) return layout; - const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { - makeBindGroupLayoutEntryBuffer(0) - }; - layout = createBindGroupLayout(device, bindGroupLayoutEntries, 1); - assert(layout); - return layout; -} - - -void WgBindGroupOpacity::releaseLayout() -{ - releaseBindGroupLayout(layout); -} - - -void WgBindGroupOpacity::initialize(WGPUDevice device, WGPUQueue queue, uint32_t uOpacity) -{ - release(); - float opacity = uOpacity / 255.0f; - uBufferOpacity = createBuffer(device, queue, &opacity, sizeof(float)); - const WGPUBindGroupEntry bindGroupEntries[] { - makeBindGroupEntryBuffer(0, uBufferOpacity) - }; - mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 1); - assert(mBindGroup); -} - - -void WgBindGroupOpacity::update(WGPUDevice device, WGPUQueue queue, uint32_t uOpacity) { - float opacity = uOpacity / 255.0f; - wgpuQueueWriteBuffer(queue, uBufferOpacity, 0, &opacity, sizeof(float)); -} - - -void WgBindGroupOpacity::release() -{ - releaseBuffer(uBufferOpacity); - releaseBindGroup(mBindGroup); -} - - WGPUBindGroupLayout WgBindGroupTexture::getLayout(WGPUDevice device) { if (layout) return layout; @@ -343,11 +301,11 @@ void WgBindGroupTexture::release() } -WGPUBindGroupLayout WgBindGroupStorageTexture::getLayout(WGPUDevice device) +WGPUBindGroupLayout WgBindGroupTextureStorage::getLayout(WGPUDevice device) { if (layout) return layout; const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { - makeBindGroupLayoutEntryStorageTexture(0) + makeBindGroupLayoutEntryStorageTexture(0, WGPUStorageTextureAccess_ReadWrite) }; layout = createBindGroupLayout(device, bindGroupLayoutEntries, 1); assert(layout); @@ -355,13 +313,13 @@ WGPUBindGroupLayout WgBindGroupStorageTexture::getLayout(WGPUDevice device) } -void WgBindGroupStorageTexture::releaseLayout() +void WgBindGroupTextureStorage::releaseLayout() { releaseBindGroupLayout(layout); } -void WgBindGroupStorageTexture::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexture) +void WgBindGroupTextureStorage::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexture) { release(); const WGPUBindGroupEntry bindGroupEntries[] { @@ -372,7 +330,7 @@ void WgBindGroupStorageTexture::initialize(WGPUDevice device, WGPUQueue queue, W } -void WgBindGroupStorageTexture::release() +void WgBindGroupTextureStorage::release() { releaseBindGroup(mBindGroup); } @@ -414,13 +372,131 @@ void WgBindGroupTextureSampled::release() releaseBindGroup(mBindGroup); } + +WGPUBindGroupLayout WgBindGroupOpacity::getLayout(WGPUDevice device) +{ + if (layout) return layout; + const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { + makeBindGroupLayoutEntryBuffer(0) + }; + layout = createBindGroupLayout(device, bindGroupLayoutEntries, 1); + assert(layout); + return layout; +} + + +void WgBindGroupOpacity::releaseLayout() +{ + releaseBindGroupLayout(layout); +} + + +void WgBindGroupOpacity::initialize(WGPUDevice device, WGPUQueue queue, uint32_t uOpacity) +{ + release(); + float opacity = uOpacity / 255.0f; + uBufferOpacity = createBuffer(device, queue, &opacity, sizeof(float)); + const WGPUBindGroupEntry bindGroupEntries[] { + makeBindGroupEntryBuffer(0, uBufferOpacity) + }; + mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 1); + assert(mBindGroup); +} + + +void WgBindGroupOpacity::release() +{ + releaseBuffer(uBufferOpacity); + releaseBindGroup(mBindGroup); +} + + +WGPUBindGroupLayout WgBindGroupBlendMethod::getLayout(WGPUDevice device) +{ + if (layout) return layout; + const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { + makeBindGroupLayoutEntryBuffer(0) + }; + layout = createBindGroupLayout(device, bindGroupLayoutEntries, 1); + assert(layout); + return layout; +} + + +void WgBindGroupBlendMethod::releaseLayout() +{ + releaseBindGroupLayout(layout); +} + + +void WgBindGroupBlendMethod::initialize(WGPUDevice device, WGPUQueue queue, BlendMethod uBlendMethod) +{ + release(); + uint32_t blendMethod = (uint32_t)uBlendMethod; + uBufferBlendMethod = createBuffer(device, queue, &blendMethod, sizeof(blendMethod)); + const WGPUBindGroupEntry bindGroupEntries[] { + makeBindGroupEntryBuffer(0, uBufferBlendMethod) + }; + mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 1); + assert(mBindGroup); +} + + +void WgBindGroupBlendMethod::release() +{ + releaseBuffer(uBufferBlendMethod); + releaseBindGroup(mBindGroup); +} + + +WGPUBindGroupLayout WgBindGroupCompositeMethod::getLayout(WGPUDevice device) +{ + if (layout) return layout; + const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { + makeBindGroupLayoutEntryBuffer(0) + }; + layout = createBindGroupLayout(device, bindGroupLayoutEntries, 1); + assert(layout); + return layout; +} + + +void WgBindGroupCompositeMethod::releaseLayout() +{ + releaseBindGroupLayout(layout); +} + + +void WgBindGroupCompositeMethod::initialize(WGPUDevice device, WGPUQueue queue, CompositeMethod uCompositeMethod) +{ + release(); + uint32_t compositeMethod = (uint32_t)uCompositeMethod; + uBufferCompositeMethod = createBuffer(device, queue, &compositeMethod, sizeof(compositeMethod)); + const WGPUBindGroupEntry bindGroupEntries[] { + makeBindGroupEntryBuffer(0, uBufferCompositeMethod) + }; + mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 1); + assert(mBindGroup); +} + + +void WgBindGroupCompositeMethod::release() +{ + releaseBuffer(uBufferCompositeMethod); + releaseBindGroup(mBindGroup); +} + //************************************************************************ // bind group pools //************************************************************************ void WgBindGroupOpacityPool::initialize(WgContext& context) { - memset(mPool, 0x00, sizeof(mPool)); + release(context); + for (uint32_t opacity = 0; opacity < 256; opacity++) { + mPool[opacity] = new WgBindGroupOpacity; + mPool[opacity]->initialize(context.device, context.queue, opacity); + } } @@ -438,11 +514,69 @@ void WgBindGroupOpacityPool::release(WgContext& context) WgBindGroupOpacity* WgBindGroupOpacityPool::allocate(WgContext& context, uint8_t opacity) { - WgBindGroupOpacity* bindGroupOpacity = mPool[opacity]; - if (!bindGroupOpacity) { - bindGroupOpacity = new WgBindGroupOpacity; - bindGroupOpacity->initialize(context.device, context.queue, opacity); - mPool[opacity] = bindGroupOpacity; - } - return bindGroupOpacity; + return mPool[opacity]; +} + + +void WgBindGroupBlendMethodPool::initialize(WgContext& context) +{ + release(context); + for (uint8_t blendMethod = (uint8_t)BlendMethod::Normal; + blendMethod <= (uint8_t)BlendMethod::SoftLight; + blendMethod++) { + mPool[blendMethod] = new WgBindGroupBlendMethod; + mPool[blendMethod]->initialize(context.device, context.queue, (BlendMethod)blendMethod); + } +} + + +void WgBindGroupBlendMethodPool::release(WgContext& context) +{ + for (uint8_t blendMethod = (uint8_t)BlendMethod::Normal; + blendMethod <= (uint8_t)BlendMethod::SoftLight; + blendMethod++) { + if (mPool[blendMethod]) { + mPool[blendMethod]->release(); + delete mPool[blendMethod]; + mPool[blendMethod] = nullptr; + } + } +} + + +WgBindGroupBlendMethod* WgBindGroupBlendMethodPool::allocate(WgContext& context, BlendMethod blendMethod) +{ + return mPool[(uint8_t)blendMethod]; +} + + +void WgBindGroupCompositeMethodPool::initialize(WgContext& context) +{ + release(context); + for (uint8_t composeMethos = (uint8_t)CompositeMethod::None; + composeMethos <= (uint8_t)CompositeMethod::DifferenceMask; + composeMethos++) { + mPool[composeMethos] = new WgBindGroupCompositeMethod; + mPool[composeMethos]->initialize(context.device, context.queue, (CompositeMethod)composeMethos); + } +} + + +void WgBindGroupCompositeMethodPool::release(WgContext& context) +{ + for (uint8_t blendMethod = (uint8_t)CompositeMethod::None; + blendMethod <= (uint8_t)CompositeMethod::DifferenceMask; + blendMethod++) { + if (mPool[blendMethod]) { + mPool[blendMethod]->release(); + delete mPool[blendMethod]; + mPool[blendMethod] = nullptr; + } + } +} + + +WgBindGroupCompositeMethod* WgBindGroupCompositeMethodPool::allocate(WgContext& context, CompositeMethod composeMethod) +{ + return mPool[(uint8_t)composeMethod]; } diff --git a/src/renderer/wg_engine/tvgWgBindGroups.h b/src/renderer/wg_engine/tvgWgBindGroups.h index 42c06c00..921bad35 100644 --- a/src/renderer/wg_engine/tvgWgBindGroups.h +++ b/src/renderer/wg_engine/tvgWgBindGroups.h @@ -106,19 +106,6 @@ struct WgBindGroupPicture : public WgBindGroup void release(); }; -// @group(1 or 2) -struct WgBindGroupOpacity : public WgBindGroup -{ - static WGPUBindGroupLayout layout; - static WGPUBindGroupLayout getLayout(WGPUDevice device); - static void releaseLayout(); - - WGPUBuffer uBufferOpacity{}; - void initialize(WGPUDevice device, WGPUQueue queue, uint32_t uOpacity); - void update(WGPUDevice device, WGPUQueue queue, uint32_t uOpacity); - void release(); -}; - // @group(0 or 1) struct WgBindGroupTexture : public WgBindGroup { @@ -132,7 +119,7 @@ struct WgBindGroupTexture : public WgBindGroup }; // @group(0 or 1) -struct WgBindGroupStorageTexture : public WgBindGroup +struct WgBindGroupTextureStorage : public WgBindGroup { static WGPUBindGroupLayout layout; static WGPUBindGroupLayout getLayout(WGPUDevice device); @@ -156,6 +143,44 @@ struct WgBindGroupTextureSampled : public WgBindGroup void release(); }; +// @group(1 or 2) +struct WgBindGroupOpacity : public WgBindGroup +{ + static WGPUBindGroupLayout layout; + static WGPUBindGroupLayout getLayout(WGPUDevice device); + static void releaseLayout(); + + WGPUBuffer uBufferOpacity{}; + void initialize(WGPUDevice device, WGPUQueue queue, uint32_t uOpacity); + void release(); +}; + +// @group(2) +struct WgBindGroupBlendMethod : public WgBindGroup +{ + static WGPUBindGroupLayout layout; + static WGPUBindGroupLayout getLayout(WGPUDevice device); + static void releaseLayout(); + + WGPUBuffer uBufferBlendMethod{}; + void initialize(WGPUDevice device, WGPUQueue queue, + BlendMethod uBlendMethod); + void release(); +}; + +// @group(2) +struct WgBindGroupCompositeMethod : public WgBindGroup +{ + static WGPUBindGroupLayout layout; + static WGPUBindGroupLayout getLayout(WGPUDevice device); + static void releaseLayout(); + + WGPUBuffer uBufferCompositeMethod{}; + void initialize(WGPUDevice device, WGPUQueue queue, + CompositeMethod uCompositeMethod); + void release(); +}; + //************************************************************************ // bind group pools //************************************************************************ @@ -163,11 +188,31 @@ struct WgBindGroupTextureSampled : public WgBindGroup class WgBindGroupOpacityPool { private: - WgBindGroupOpacity* mPool[256]; + WgBindGroupOpacity* mPool[256]{}; public: void initialize(WgContext& context); void release(WgContext& context); WgBindGroupOpacity* allocate(WgContext& context, uint8_t opacity); }; +class WgBindGroupBlendMethodPool +{ +private: + WgBindGroupBlendMethod* mPool[(uint8_t)BlendMethod::SoftLight + 1]{}; +public: + void initialize(WgContext& context); + void release(WgContext& context); + WgBindGroupBlendMethod* allocate(WgContext& context, BlendMethod blendMethod); +}; + +class WgBindGroupCompositeMethodPool +{ +private: + WgBindGroupCompositeMethod* mPool[(uint8_t)CompositeMethod::DifferenceMask + 1]{}; +public: + void initialize(WgContext& context); + void release(WgContext& context); + WgBindGroupCompositeMethod* allocate(WgContext& context, CompositeMethod composeMethod); +}; + #endif // _TVG_WG_BIND_GROUPS_H_ diff --git a/src/renderer/wg_engine/tvgWgCommon.cpp b/src/renderer/wg_engine/tvgWgCommon.cpp index e0d6e8d2..4c5bb430 100644 --- a/src/renderer/wg_engine/tvgWgCommon.cpp +++ b/src/renderer/wg_engine/tvgWgCommon.cpp @@ -88,11 +88,19 @@ void WgContext::initialize() queue = wgpuDeviceGetQueue(device); assert(queue); + + // create default nearest and linear samplers + samplerNearest = createSampler(WGPUFilterMode_Nearest, WGPUMipmapFilterMode_Nearest); + assert(samplerNearest); + samplerLinear = createSampler(WGPUFilterMode_Linear, WGPUMipmapFilterMode_Linear); + assert(samplerLinear); } void WgContext::release() { + releaseSampler(samplerNearest); + releaseSampler(samplerLinear); if (device) { wgpuDeviceDestroy(device); wgpuDeviceRelease(device); @@ -115,7 +123,26 @@ void WgContext::executeCommandEncoder(WGPUCommandEncoder commandEncoder) } -WGPUTexture WgContext::createTexture2d(WGPUTextureUsage usage, WGPUTextureFormat format, uint32_t width, uint32_t height, char const * label) { +WGPUSampler WgContext::createSampler(WGPUFilterMode minFilter, WGPUMipmapFilterMode mipmapFilter) +{ + WGPUSamplerDescriptor samplerDesc{}; + samplerDesc.nextInChain = nullptr; + samplerDesc.label = "The sampler"; + samplerDesc.addressModeU = WGPUAddressMode_ClampToEdge; + samplerDesc.addressModeV = WGPUAddressMode_ClampToEdge; + samplerDesc.addressModeW = WGPUAddressMode_ClampToEdge; + samplerDesc.magFilter = minFilter; + samplerDesc.minFilter = minFilter; + samplerDesc.mipmapFilter = mipmapFilter; + samplerDesc.lodMinClamp = 0.0f; + samplerDesc.lodMaxClamp = 32.0f; + samplerDesc.compare = WGPUCompareFunction_Undefined; + samplerDesc.maxAnisotropy = 1; + return wgpuDeviceCreateSampler(device, &samplerDesc); +} + + +WGPUTexture WgContext::createTexture2d(WGPUTextureUsageFlags usage, WGPUTextureFormat format, uint32_t width, uint32_t height, char const * label) { WGPUTextureDescriptor textureDesc{}; textureDesc.nextInChain = nullptr; textureDesc.label = label; @@ -147,7 +174,16 @@ WGPUTextureView WgContext::createTextureView2d(WGPUTexture texture, WGPU_NULLABL }; -void WgContext::releaseTexture(WGPUTexture& texture) { +void WgContext::releaseSampler(WGPUSampler& sampler) +{ + if (sampler) { + wgpuSamplerRelease(sampler); + sampler = nullptr; + } +} + +void WgContext::releaseTexture(WGPUTexture& texture) +{ if (texture) { wgpuTextureDestroy(texture); wgpuTextureRelease(texture); @@ -157,7 +193,8 @@ void WgContext::releaseTexture(WGPUTexture& texture) { } -void WgContext::releaseTextureView(WGPUTextureView& textureView) { +void WgContext::releaseTextureView(WGPUTextureView& textureView) +{ if (textureView) wgpuTextureViewRelease(textureView); textureView = nullptr; } @@ -226,7 +263,7 @@ WGPUBindGroupLayoutEntry WgBindGroup::makeBindGroupLayoutEntryBuffer(uint32_t bi WGPUBindGroupLayoutEntry bindGroupLayoutEntry{}; bindGroupLayoutEntry.nextInChain = nullptr; bindGroupLayoutEntry.binding = binding; - bindGroupLayoutEntry.visibility = WGPUShaderStage_Vertex | WGPUShaderStage_Fragment; + bindGroupLayoutEntry.visibility = WGPUShaderStage_Vertex | WGPUShaderStage_Fragment | WGPUShaderStage_Compute; bindGroupLayoutEntry.buffer.nextInChain = nullptr; bindGroupLayoutEntry.buffer.type = WGPUBufferBindingType_Uniform; bindGroupLayoutEntry.buffer.hasDynamicOffset = false; @@ -261,14 +298,14 @@ WGPUBindGroupLayoutEntry WgBindGroup::makeBindGroupLayoutEntryTexture(uint32_t b } -WGPUBindGroupLayoutEntry WgBindGroup::makeBindGroupLayoutEntryStorageTexture(uint32_t binding) +WGPUBindGroupLayoutEntry WgBindGroup::makeBindGroupLayoutEntryStorageTexture(uint32_t binding, WGPUStorageTextureAccess access) { WGPUBindGroupLayoutEntry bindGroupLayoutEntry{}; bindGroupLayoutEntry.nextInChain = nullptr; bindGroupLayoutEntry.binding = binding; bindGroupLayoutEntry.visibility = WGPUShaderStage_Fragment | WGPUShaderStage_Compute; bindGroupLayoutEntry.storageTexture.nextInChain = nullptr; - bindGroupLayoutEntry.storageTexture.access = WGPUStorageTextureAccess_ReadWrite; + bindGroupLayoutEntry.storageTexture.access = access; bindGroupLayoutEntry.storageTexture.format = WGPUTextureFormat_RGBA8Unorm; bindGroupLayoutEntry.storageTexture.viewDimension = WGPUTextureViewDimension_2D; return bindGroupLayoutEntry; @@ -427,11 +464,11 @@ WGPUBlendState WgRenderPipeline::makeBlendState() { WGPUBlendState blendState{}; blendState.color.operation = WGPUBlendOperation_Add; - blendState.color.srcFactor = WGPUBlendFactor_SrcAlpha; - blendState.color.dstFactor = WGPUBlendFactor_OneMinusSrcAlpha; + blendState.color.srcFactor = WGPUBlendFactor_One; + blendState.color.dstFactor = WGPUBlendFactor_Zero; blendState.alpha.operation = WGPUBlendOperation_Add; blendState.alpha.srcFactor = WGPUBlendFactor_One; - blendState.alpha.dstFactor = WGPUBlendFactor_One; + blendState.alpha.dstFactor = WGPUBlendFactor_Zero; return blendState; } diff --git a/src/renderer/wg_engine/tvgWgCommon.h b/src/renderer/wg_engine/tvgWgCommon.h index 5528eaa1..ce36b5ea 100644 --- a/src/renderer/wg_engine/tvgWgCommon.h +++ b/src/renderer/wg_engine/tvgWgCommon.h @@ -40,6 +40,9 @@ struct WgContext { WGPUAdapterProperties adapterProperties{}; WGPUSupportedLimits supportedLimits{}; + WGPUSampler samplerNearest{}; + WGPUSampler samplerLinear{}; + WgPipelines* pipelines{}; // external handle (do not release) void initialize(); @@ -47,8 +50,10 @@ struct WgContext { void executeCommandEncoder(WGPUCommandEncoder commandEncoder); - WGPUTexture createTexture2d(WGPUTextureUsage usage, WGPUTextureFormat format, uint32_t width, uint32_t height, char const * label); + WGPUSampler createSampler(WGPUFilterMode minFilter, WGPUMipmapFilterMode mipmapFilter); + WGPUTexture createTexture2d(WGPUTextureUsageFlags usage, WGPUTextureFormat format, uint32_t width, uint32_t height, char const * label); WGPUTextureView createTextureView2d(WGPUTexture texture, WGPU_NULLABLE char const * label); + void releaseSampler(WGPUSampler& sampler); void releaseTexture(WGPUTexture& texture); void releaseTextureView(WGPUTextureView& textureView); }; @@ -67,7 +72,7 @@ struct WgBindGroup static WGPUBindGroupLayoutEntry makeBindGroupLayoutEntryBuffer(uint32_t binding); static WGPUBindGroupLayoutEntry makeBindGroupLayoutEntrySampler(uint32_t binding); static WGPUBindGroupLayoutEntry makeBindGroupLayoutEntryTexture(uint32_t binding); - static WGPUBindGroupLayoutEntry makeBindGroupLayoutEntryStorageTexture(uint32_t binding); + static WGPUBindGroupLayoutEntry makeBindGroupLayoutEntryStorageTexture(uint32_t binding, WGPUStorageTextureAccess access); static WGPUBuffer createBuffer(WGPUDevice device, WGPUQueue queue, const void *data, size_t size); static WGPUBindGroup createBindGroup(WGPUDevice device, WGPUBindGroupLayout layout, const WGPUBindGroupEntry* bindGroupEntries, uint32_t count); @@ -123,6 +128,9 @@ public: static void destroyRenderPipeline(WGPURenderPipeline& renderPipeline); }; +#define WG_COMPUTE_WORKGROUP_SIZE_X 8 +#define WG_COMPUTE_WORKGROUP_SIZE_Y 8 + struct WgComputePipeline: public WgPipeline { protected: diff --git a/src/renderer/wg_engine/tvgWgPipelines.cpp b/src/renderer/wg_engine/tvgWgPipelines.cpp index 099b7319..209a3692 100644 --- a/src/renderer/wg_engine/tvgWgPipelines.cpp +++ b/src/renderer/wg_engine/tvgWgPipelines.cpp @@ -25,6 +25,10 @@ #define ARRAY_ELEMENTS_COUNT(arr) sizeof(arr)/sizeof(arr[0]) +//************************************************************************ +// graphics pipelines +//************************************************************************ + void WgPipelineFillShape::initialize(WGPUDevice device) { // vertex attributes settings @@ -222,120 +226,64 @@ void WgPipelineImage::initialize(WGPUDevice device) shaderSource, shaderLabel, pipelineLabel); } +//************************************************************************ +// compute pipelines +//************************************************************************ -void WgPipelineBlit::initialize(WGPUDevice device) +void WgPipelineClear::initialize(WGPUDevice device) { - // vertex and buffers settings - WGPUVertexAttribute vertexAttributesPos = { WGPUVertexFormat_Float32x2, sizeof(float) * 0, 0 }; - WGPUVertexAttribute vertexAttributesTex = { WGPUVertexFormat_Float32x2, sizeof(float) * 0, 1 }; - WGPUVertexBufferLayout vertexBufferLayouts[] = { - makeVertexBufferLayout(&vertexAttributesPos, 1, sizeof(float) * 2), - makeVertexBufferLayout(&vertexAttributesTex, 1, sizeof(float) * 2) - }; - // bind groups and layouts WGPUBindGroupLayout bindGroupLayouts[] = { - WgBindGroupTextureSampled::getLayout(device), + WgBindGroupTextureStorage::getLayout(device) + }; + + // sheder source and labels + auto shaderSource = cShaderSource_PipelineComputeClear; + auto shaderLabel = "The compute shader clear"; + auto pipelineLabel = "The compute pipeline clear"; + + // allocate all pipeline handles + allocate(device, + bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts), + shaderSource, shaderLabel, pipelineLabel); +} + + +void WgPipelineBlend::initialize(WGPUDevice device) +{ + // bind groups and layouts + WGPUBindGroupLayout bindGroupLayouts[] = { + WgBindGroupTextureStorage::getLayout(device), + WgBindGroupTextureStorage::getLayout(device), + WgBindGroupBlendMethod::getLayout(device) + }; + + // sheder source and labels + auto shaderSource = cShaderSource_PipelineComputeBlend; + auto shaderLabel = "The compute shader blend"; + auto pipelineLabel = "The compute pipeline blend"; + + // 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[] = { + WgBindGroupTextureStorage::getLayout(device), + WgBindGroupTextureStorage::getLayout(device), + WgBindGroupCompositeMethod::getLayout(device), WgBindGroupOpacity::getLayout(device) }; - // stencil function - WGPUCompareFunction stencilFuncion = WGPUCompareFunction_Always; - WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero; - // sheder source and labels - auto shaderSource = cShaderSource_PipelineBlit; - auto shaderLabel = "The shader blit"; - auto pipelineLabel = "The render pipeline blit"; - - // allocate all pipeline handles - allocate(device, - vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts), - bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts), - stencilFuncion, stencilOperation, - shaderSource, shaderLabel, pipelineLabel); -} - - -void WgPipelineBlitColor::initialize(WGPUDevice device) -{ - // vertex and buffers settings - WGPUVertexAttribute vertexAttributesPos = { WGPUVertexFormat_Float32x2, sizeof(float) * 0, 0 }; - WGPUVertexAttribute vertexAttributesTex = { WGPUVertexFormat_Float32x2, sizeof(float) * 0, 1 }; - WGPUVertexBufferLayout vertexBufferLayouts[] = { - makeVertexBufferLayout(&vertexAttributesPos, 1, sizeof(float) * 2), - makeVertexBufferLayout(&vertexAttributesTex, 1, sizeof(float) * 2) - }; - - // bind groups and layouts - WGPUBindGroupLayout bindGroupLayouts[] = { - WgBindGroupTextureSampled::getLayout(device) - }; - - // stencil function - WGPUCompareFunction stencilFuncion = WGPUCompareFunction_Always; - WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero; - - // sheder source and labels - auto shaderSource = cShaderSource_PipelineBlitColor; - auto shaderLabel = "The shader blit color"; - auto pipelineLabel = "The render pipeline blit color"; - - // allocate all pipeline handles - allocate(device, - vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts), - bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts), - stencilFuncion, stencilOperation, - shaderSource, shaderLabel, pipelineLabel); -} - - -void WgPipelineComposition::initialize(WGPUDevice device, const char* shaderSrc) -{ - // vertex and buffers settings - WGPUVertexAttribute vertexAttributesPos = { WGPUVertexFormat_Float32x2, sizeof(float) * 0, 0 }; - WGPUVertexAttribute vertexAttributesTex = { WGPUVertexFormat_Float32x2, sizeof(float) * 0, 1 }; - WGPUVertexBufferLayout vertexBufferLayouts[] = { - makeVertexBufferLayout(&vertexAttributesPos, 1, sizeof(float) * 2), - makeVertexBufferLayout(&vertexAttributesTex, 1, sizeof(float) * 2) - }; - - // bind groups and layouts - WGPUBindGroupLayout bindGroupLayouts[] = { - WgBindGroupTextureSampled::getLayout(device), - WgBindGroupTextureSampled::getLayout(device) - }; - - // stencil function - WGPUCompareFunction stencilFuncion = WGPUCompareFunction_Always; - WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero; - - // sheder source and labels - auto shaderSource = shaderSrc; - auto shaderLabel = "The shader compose alpha mask"; - auto pipelineLabel = "The render pipeline compose alpha mask"; - - // allocate all pipeline handles - allocate(device, - vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts), - bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts), - stencilFuncion, stencilOperation, - shaderSource, shaderLabel, pipelineLabel); -} - -void WgPipelineBlend::initialize(WGPUDevice device, const char* shaderSrc) -{ - // bind groups and layouts - WGPUBindGroupLayout bindGroupLayouts[] = { - //WgBindGroupTexture::getLayout(device), - WgBindGroupStorageTexture::getLayout(device), - WgBindGroupStorageTexture::getLayout(device) - }; - - // sheder source and labels - auto shaderSource = shaderSrc; - auto shaderLabel = "The compute shader blend"; - auto pipelineLabel = "The compute pipeline blend"; + auto shaderSource = cShaderSource_PipelineComputeCompose; + auto shaderLabel = "The compute shader compose"; + auto pipelineLabel = "The compute pipeline compose"; // allocate all pipeline handles allocate(device, @@ -356,20 +304,10 @@ void WgPipelines::initialize(WgContext& context) linear.initialize(context.device); radial.initialize(context.device); image.initialize(context.device); - // blit pipelines - blit.initialize(context.device); - blitColor.initialize(context.device); - // composition pipelines - compAlphaMask.initialize(context.device, cShaderSource_PipelineCompAlphaMask); - compInvAlphaMask.initialize(context.device, cShaderSource_PipelineCompInvAlphaMask); - compLumaMask.initialize(context.device, cShaderSource_PipelineCompLumaMask); - compInvLumaMask.initialize(context.device, cShaderSource_PipelineCompInvLumaMask); - compAddMask.initialize(context.device, cShaderSource_PipelineCompAddMask); - compSubtractMask.initialize(context.device, cShaderSource_PipelineCompSubtractMask); - compIntersectMask.initialize(context.device, cShaderSource_PipelineCompIntersectMask); - compDifferenceMask.initialize(context.device, cShaderSource_PipelineCompDifferenceMask); // compute pipelines - computeBlend.initialize(context.device, cShaderSource_PipelineComputeBlend); + computeClear.initialize(context.device); + computeBlend.initialize(context.device); + computeCompose.initialize(context.device); // store pipelines to context context.pipelines = this; } @@ -378,7 +316,7 @@ void WgPipelines::initialize(WgContext& context) void WgPipelines::release() { WgBindGroupTextureSampled::releaseLayout(); - WgBindGroupStorageTexture::releaseLayout(); + WgBindGroupTextureStorage::releaseLayout(); WgBindGroupTexture::releaseLayout(); WgBindGroupOpacity::releaseLayout(); WgBindGroupPicture::releaseLayout(); @@ -388,19 +326,9 @@ void WgPipelines::release() WgBindGroupPaint::releaseLayout(); WgBindGroupCanvas::releaseLayout(); // compute pipelines + computeCompose.release(); computeBlend.release(); - // composition pipelines - compDifferenceMask.release(); - compIntersectMask.release(); - compSubtractMask.release(); - compAddMask.release(); - compInvLumaMask.release(); - compLumaMask.release(); - compInvAlphaMask.release(); - compAlphaMask.release(); - // blit pipelines - blitColor.release(); - blit.release(); + computeClear.release(); // fill pipelines image.release(); radial.release(); @@ -409,21 +337,3 @@ void WgPipelines::release() fillStroke.release(); fillShape.release(); } - - -WgPipelineComposition* WgPipelines::getCompositionPipeline(CompositeMethod method) -{ - switch (method) { - case CompositeMethod::ClipPath: - case CompositeMethod::AlphaMask: return &compAlphaMask; break; - case CompositeMethod::InvAlphaMask: return &compInvAlphaMask; break; - case CompositeMethod::LumaMask: return &compLumaMask; break; - case CompositeMethod::InvLumaMask: return &compInvLumaMask; break; - case CompositeMethod::AddMask: return &compAddMask; break; - case CompositeMethod::SubtractMask: return &compSubtractMask; break; - case CompositeMethod::IntersectMask: return &compIntersectMask; break; - case CompositeMethod::DifferenceMask: return &compDifferenceMask; break; - default: return nullptr; break; - } - return nullptr; -} diff --git a/src/renderer/wg_engine/tvgWgPipelines.h b/src/renderer/wg_engine/tvgWgPipelines.h index 8c34b23c..99097828 100644 --- a/src/renderer/wg_engine/tvgWgPipelines.h +++ b/src/renderer/wg_engine/tvgWgPipelines.h @@ -99,59 +99,52 @@ struct WgPipelineImage: public WgRenderPipeline } }; -struct WgPipelineBlit: public WgRenderPipeline -{ - void initialize(WGPUDevice device) override; - void use(WGPURenderPassEncoder encoder, - WgBindGroupTextureSampled& groupTexSampled, - WgBindGroupOpacity& groupOpacity) - { - set(encoder); - groupTexSampled.set(encoder, 0); - groupOpacity.set(encoder, 1); - } -}; - -struct WgPipelineBlitColor: public WgRenderPipeline -{ - void initialize(WGPUDevice device) override; - void use(WGPURenderPassEncoder encoder, - WgBindGroupTextureSampled& groupTexSampled) - { - set(encoder); - groupTexSampled.set(encoder, 0); - } -}; - -struct WgPipelineComposition: public WgRenderPipeline -{ - void initialize(WGPUDevice device) override {}; - void initialize(WGPUDevice device, const char* shaderSrc); - void use(WGPURenderPassEncoder encoder, - WgBindGroupTextureSampled& groupTexSampledSrc, - WgBindGroupTextureSampled& groupTexSampledMsk) - { - set(encoder); - groupTexSampledSrc.set(encoder, 0); - groupTexSampledMsk.set(encoder, 1); - } -}; - //***************************************************************************** // compute pipelines //***************************************************************************** +struct WgPipelineClear: public WgComputePipeline +{ + void initialize(WGPUDevice device) override; + void use(WGPUComputePassEncoder encoder, + WgBindGroupTextureStorage& groupTexDst) + { + set(encoder); + groupTexDst.set(encoder, 0); + } +}; + + struct WgPipelineBlend: public WgComputePipeline { - void initialize(WGPUDevice device) override {}; - void initialize(WGPUDevice device, const char* shaderSrc); + void initialize(WGPUDevice device) override; void use(WGPUComputePassEncoder encoder, - WgBindGroupStorageTexture& groupTexSrc, - WgBindGroupStorageTexture& groupTexDst) + WgBindGroupTextureStorage& groupTexSrc, + WgBindGroupTextureStorage& groupTexDst, + WgBindGroupBlendMethod& blendMethod) { set(encoder); groupTexSrc.set(encoder, 0); groupTexDst.set(encoder, 1); + blendMethod.set(encoder, 2); + } +}; + + +struct WgPipelineCompose: public WgComputePipeline +{ + void initialize(WGPUDevice device) override; + void use(WGPUComputePassEncoder encoder, + WgBindGroupTextureStorage& groupTexSrc, + WgBindGroupTextureStorage& groupTexMsk, + WgBindGroupCompositeMethod& groupComposeMethod, + WgBindGroupOpacity& groupOpacity) + { + set(encoder); + groupTexSrc.set(encoder, 0); + groupTexMsk.set(encoder, 1); + groupComposeMethod.set(encoder, 2); + groupOpacity.set(encoder, 3); } }; @@ -161,30 +154,20 @@ struct WgPipelineBlend: public WgComputePipeline struct WgPipelines { + // render pipelines WgPipelineFillShape fillShape; WgPipelineFillStroke fillStroke; WgPipelineSolid solid; WgPipelineLinear linear; WgPipelineRadial radial; WgPipelineImage image; - WgPipelineBlit blit; - WgPipelineBlitColor blitColor; - // composition pipelines - WgPipelineComposition compAlphaMask; - WgPipelineComposition compInvAlphaMask; - WgPipelineComposition compLumaMask; - WgPipelineComposition compInvLumaMask; - WgPipelineComposition compAddMask; - WgPipelineComposition compSubtractMask; - WgPipelineComposition compIntersectMask; - WgPipelineComposition compDifferenceMask; // compute pipelines + WgPipelineClear computeClear; WgPipelineBlend computeBlend; + WgPipelineCompose computeCompose; void initialize(WgContext& context); void release(); - - WgPipelineComposition* getCompositionPipeline(CompositeMethod method); }; #endif // _TVG_WG_PIPELINES_H_ diff --git a/src/renderer/wg_engine/tvgWgRenderData.cpp b/src/renderer/wg_engine/tvgWgRenderData.cpp index 30169c5f..bceb646c 100644 --- a/src/renderer/wg_engine/tvgWgRenderData.cpp +++ b/src/renderer/wg_engine/tvgWgRenderData.cpp @@ -144,48 +144,12 @@ void WgImageData::update(WgContext& context, Surface* surface) { release(context); assert(surface); - // sampler descriptor - WGPUSamplerDescriptor samplerDesc{}; - samplerDesc.nextInChain = nullptr; - samplerDesc.label = "The shape sampler"; - samplerDesc.addressModeU = WGPUAddressMode_ClampToEdge; - samplerDesc.addressModeV = WGPUAddressMode_ClampToEdge; - samplerDesc.addressModeW = WGPUAddressMode_ClampToEdge; - samplerDesc.magFilter = WGPUFilterMode_Nearest; - samplerDesc.minFilter = WGPUFilterMode_Nearest; - samplerDesc.mipmapFilter = WGPUMipmapFilterMode_Nearest; - samplerDesc.lodMinClamp = 0.0f; - samplerDesc.lodMaxClamp = 32.0f; - samplerDesc.compare = WGPUCompareFunction_Undefined; - samplerDesc.maxAnisotropy = 1; - sampler = wgpuDeviceCreateSampler(context.device, &samplerDesc); - assert(sampler); - // texture descriptor - WGPUTextureDescriptor textureDesc{}; - textureDesc.nextInChain = nullptr; - textureDesc.label = "The shape texture"; - textureDesc.usage = WGPUTextureUsage_TextureBinding | WGPUTextureUsage_CopyDst; - textureDesc.dimension = WGPUTextureDimension_2D; - textureDesc.size = { surface->w, surface->h, 1 }; - textureDesc.format = WGPUTextureFormat_RGBA8Unorm; - textureDesc.mipLevelCount = 1; - textureDesc.sampleCount = 1; - textureDesc.viewFormatCount = 0; - textureDesc.viewFormats = nullptr; - texture = wgpuDeviceCreateTexture(context.device, &textureDesc); + texture = context.createTexture2d( + WGPUTextureUsage_TextureBinding | WGPUTextureUsage_CopyDst, + WGPUTextureFormat_RGBA8Unorm, + surface->w, surface->h, "The shape texture"); assert(texture); - // texture view descriptor - WGPUTextureViewDescriptor textureViewDesc{}; - textureViewDesc.nextInChain = nullptr; - textureViewDesc.label = "The shape texture view"; - textureViewDesc.format = WGPUTextureFormat_RGBA8Unorm; - textureViewDesc.dimension = WGPUTextureViewDimension_2D; - textureViewDesc.baseMipLevel = 0; - textureViewDesc.mipLevelCount = 1; - textureViewDesc.baseArrayLayer = 0; - textureViewDesc.arrayLayerCount = 1; - textureViewDesc.aspect = WGPUTextureAspect_All; - textureView = wgpuTextureCreateView(texture, &textureViewDesc); + textureView = context.createTextureView2d(texture, "The shape texture view"); assert(textureView); // update texture data WGPUImageCopyTexture imageCopyTexture{}; @@ -209,19 +173,8 @@ void WgImageData::update(WgContext& context, Surface* surface) void WgImageData::release(WgContext& context) { - if (textureView) { - wgpuTextureViewRelease(textureView); - textureView = nullptr; - } - if (texture) { - wgpuTextureDestroy(texture); - wgpuTextureRelease(texture); - texture = nullptr; - } - if (sampler) { - wgpuSamplerRelease(sampler); - sampler = nullptr; - } + context.releaseTextureView(textureView); + context.releaseTexture(texture); }; //*********************************************************************** diff --git a/src/renderer/wg_engine/tvgWgRenderData.h b/src/renderer/wg_engine/tvgWgRenderData.h index 55cf964b..1e99036e 100644 --- a/src/renderer/wg_engine/tvgWgRenderData.h +++ b/src/renderer/wg_engine/tvgWgRenderData.h @@ -45,7 +45,6 @@ struct WgMeshDataGroup { }; struct WgImageData { - WGPUSampler sampler{}; WGPUTexture texture{}; WGPUTextureView textureView{}; diff --git a/src/renderer/wg_engine/tvgWgRenderTarget.cpp b/src/renderer/wg_engine/tvgWgRenderTarget.cpp index 193438ea..11c9287f 100644 --- a/src/renderer/wg_engine/tvgWgRenderTarget.cpp +++ b/src/renderer/wg_engine/tvgWgRenderTarget.cpp @@ -22,133 +22,128 @@ #include "tvgWgRenderTarget.h" +//***************************************************************************** +// render target +//***************************************************************************** + void WgRenderTarget::initialize(WgContext& context, uint32_t w, uint32_t h) { release(context); - // sampler descriptor - WGPUSamplerDescriptor samplerDesc{}; - samplerDesc.nextInChain = nullptr; - samplerDesc.label = "The target sampler"; - samplerDesc.addressModeU = WGPUAddressMode_ClampToEdge; - samplerDesc.addressModeV = WGPUAddressMode_ClampToEdge; - samplerDesc.addressModeW = WGPUAddressMode_ClampToEdge; - samplerDesc.magFilter = WGPUFilterMode_Nearest; - samplerDesc.minFilter = WGPUFilterMode_Nearest; - samplerDesc.mipmapFilter = WGPUMipmapFilterMode_Nearest; - samplerDesc.lodMinClamp = 0.0f; - samplerDesc.lodMaxClamp = 32.0f; - samplerDesc.compare = WGPUCompareFunction_Undefined; - samplerDesc.maxAnisotropy = 1; - sampler = wgpuDeviceCreateSampler(context.device, &samplerDesc); - assert(sampler); - // texture descriptor - WGPUTextureDescriptor textureDescColor{}; - textureDescColor.nextInChain = nullptr; - textureDescColor.label = "The target texture color"; - textureDescColor.usage = WGPUTextureUsage_RenderAttachment | WGPUTextureUsage_TextureBinding | WGPUTextureUsage_CopyDst | WGPUTextureUsage_StorageBinding; - textureDescColor.dimension = WGPUTextureDimension_2D; - textureDescColor.size = { w, h, 1 }; - textureDescColor.format = WGPUTextureFormat_RGBA8Unorm; - textureDescColor.mipLevelCount = 1; - textureDescColor.sampleCount = 1; - textureDescColor.viewFormatCount = 0; - textureDescColor.viewFormats = nullptr; - mTextureColor = wgpuDeviceCreateTexture(context.device, &textureDescColor); - assert(mTextureColor); - // texture view descriptor - WGPUTextureViewDescriptor textureViewDescColor{}; - textureViewDescColor.nextInChain = nullptr; - textureViewDescColor.label = "The target texture view color"; - textureViewDescColor.format = WGPUTextureFormat_RGBA8Unorm; - textureViewDescColor.dimension = WGPUTextureViewDimension_2D; - textureViewDescColor.baseMipLevel = 0; - textureViewDescColor.mipLevelCount = 1; - textureViewDescColor.baseArrayLayer = 0; - textureViewDescColor.arrayLayerCount = 1; - textureViewDescColor.aspect = WGPUTextureAspect_All; - textureViewColor = wgpuTextureCreateView(mTextureColor, &textureViewDescColor); - assert(textureViewColor); - // stencil texture - WGPUTextureDescriptor textureDescStencil{}; - textureDescStencil.nextInChain = nullptr; - textureDescStencil.label = "The target texture stencil"; - textureDescStencil.usage = WGPUTextureUsage_RenderAttachment; - textureDescStencil.dimension = WGPUTextureDimension_2D; - textureDescStencil.size = { w, h, 1 }; - textureDescStencil.format = WGPUTextureFormat_Stencil8; - textureDescStencil.mipLevelCount = 1; - textureDescStencil.sampleCount = 1; - textureDescStencil.viewFormatCount = 0; - textureDescStencil.viewFormats = nullptr; - mTextureStencil = wgpuDeviceCreateTexture(context.device, &textureDescStencil); - assert(mTextureStencil); - // texture view descriptor - WGPUTextureViewDescriptor textureViewDescStencil{}; - textureViewDescStencil.nextInChain = nullptr; - textureViewDescStencil.label = "The target texture view stncil"; - textureViewDescStencil.format = WGPUTextureFormat_Stencil8; - textureViewDescStencil.dimension = WGPUTextureViewDimension_2D; - textureViewDescStencil.baseMipLevel = 0; - textureViewDescStencil.mipLevelCount = 1; - textureViewDescStencil.baseArrayLayer = 0; - textureViewDescStencil.arrayLayerCount = 1; - textureViewDescStencil.aspect = WGPUTextureAspect_All; - textureViewStencil = wgpuTextureCreateView(mTextureStencil, &textureViewDescStencil); - assert(textureViewStencil); + // create color and stencil textures + texColor = context.createTexture2d( + WGPUTextureUsage_RenderAttachment | WGPUTextureUsage_StorageBinding, + WGPUTextureFormat_RGBA8Unorm, + w, h, "The target texture color"); + texStencil = context.createTexture2d( + WGPUTextureUsage_RenderAttachment, + WGPUTextureFormat_Stencil8, + w, h, "The target texture stencil"); + assert(texColor); + assert(texStencil); + texViewColor = context.createTextureView2d(texColor, "The target texture view color"); + texViewStencil = context.createTextureView2d(texStencil, "The target texture view stencil"); + assert(texViewColor); + assert(texViewStencil); // initialize bind group for blitting - bindGroupTex.initialize(context.device, context.queue, textureViewColor); - bindGroupStorageTex.initialize(context.device, context.queue, textureViewColor); - bindGroupTexSampled.initialize(context.device, context.queue, sampler, textureViewColor); + bindGroupTexStorage.initialize(context.device, context.queue, texViewColor); // initialize window binding groups WgShaderTypeMat4x4f viewMat(w, h); - mBindGroupCanvasWnd.initialize(context.device, context.queue, viewMat); - WgGeometryData geometryDataWnd; - geometryDataWnd.appendBlitBox(); - mMeshDataCanvasWnd.update(context, &geometryDataWnd); + mBindGroupCanvas.initialize(context.device, context.queue, viewMat); mPipelines = context.pipelines; } void WgRenderTarget::release(WgContext& context) { - mMeshDataCanvasWnd.release(context); - mBindGroupCanvasWnd.release(); - bindGroupTexSampled.release(); - bindGroupStorageTex.release(); - bindGroupTex.release(); - if (mTextureStencil) { - wgpuTextureDestroy(mTextureStencil); - wgpuTextureRelease(mTextureStencil); - mTextureStencil = nullptr; - } - if (textureViewStencil) wgpuTextureViewRelease(textureViewStencil); - textureViewStencil = nullptr; - if (mTextureColor) { - wgpuTextureDestroy(mTextureColor); - wgpuTextureRelease(mTextureColor); - mTextureColor = nullptr; - } - if (textureViewColor) wgpuTextureViewRelease(textureViewColor); - textureViewColor = nullptr; - if (sampler) wgpuSamplerRelease(sampler); - sampler = nullptr; + mBindGroupCanvas.release(); + bindGroupTexStorage.release(); + context.releaseTextureView(texViewStencil); + context.releaseTextureView(texViewColor); + context.releaseTexture(texStencil); + context.releaseTexture(texColor); } -void WgRenderTarget::beginRenderPass(WGPUCommandEncoder commandEncoder, bool clear) +void WgRenderTarget::renderShape(WGPUCommandEncoder commandEncoder, WgRenderDataShape* renderData) { - beginRenderPass(commandEncoder, textureViewColor, clear); + assert(renderData); + assert(commandEncoder); + WGPURenderPassEncoder renderPassEncoder = beginRenderPass(commandEncoder); + drawShape(renderPassEncoder, renderData); + drawStroke(renderPassEncoder, renderData); + endRenderPass(renderPassEncoder); } -void WgRenderTarget::beginRenderPass(WGPUCommandEncoder commandEncoder, WGPUTextureView colorAttachement, bool clear) +void WgRenderTarget::renderPicture(WGPUCommandEncoder commandEncoder, WgRenderDataPicture* renderData) { - assert(!mRenderPassEncoder); + assert(renderData); + assert(commandEncoder); + WGPURenderPassEncoder renderPassEncoder = beginRenderPass(commandEncoder); + wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 0); + mPipelines->image.use(renderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint, renderData->bindGroupPicture); + renderData->meshData.drawImage(renderPassEncoder); + endRenderPass(renderPassEncoder); +} + + +void WgRenderTarget::drawShape(WGPURenderPassEncoder renderPassEncoder, WgRenderDataShape* renderData) +{ + assert(renderData); + assert(renderPassEncoder); + // draw shape geometry + wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 0); + for (uint32_t i = 0; i < renderData->meshGroupShapes.meshes.count; i++) { + // draw to stencil (first pass) + mPipelines->fillShape.use(renderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint); + renderData->meshGroupShapes.meshes[i]->draw(renderPassEncoder); + // fill shape (second pass) + WgRenderSettings& settings = renderData->renderSettingsShape; + if (settings.fillType == WgRenderSettingsType::Solid) + mPipelines->solid.use(renderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint, settings.bindGroupSolid); + else if (settings.fillType == WgRenderSettingsType::Linear) + mPipelines->linear.use(renderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint, settings.bindGroupLinear); + else if (settings.fillType == WgRenderSettingsType::Radial) + mPipelines->radial.use(renderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint, settings.bindGroupRadial); + renderData->meshBBoxShapes.draw(renderPassEncoder); + } +} + + +void WgRenderTarget::drawStroke(WGPURenderPassEncoder renderPassEncoder, WgRenderDataShape* renderData) +{ + assert(renderData); + assert(renderPassEncoder); + // draw stroke geometry + if (renderData->meshGroupStrokes.meshes.count > 0) { + // draw strokes to stencil (first pass) + wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 255); + for (uint32_t i = 0; i < renderData->meshGroupStrokes.meshes.count; i++) { + mPipelines->fillStroke.use(renderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint); + renderData->meshGroupStrokes.meshes[i]->draw(renderPassEncoder); + } + // fill shape (second pass) + wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 0); + WgRenderSettings& settings = renderData->renderSettingsStroke; + if (settings.fillType == WgRenderSettingsType::Solid) + mPipelines->solid.use(renderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint, settings.bindGroupSolid); + else if (settings.fillType == WgRenderSettingsType::Linear) + mPipelines->linear.use(renderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint, settings.bindGroupLinear); + else if (settings.fillType == WgRenderSettingsType::Radial) + mPipelines->radial.use(renderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint, settings.bindGroupRadial); + renderData->meshBBoxStrokes.draw(renderPassEncoder); + } +} + + +WGPURenderPassEncoder WgRenderTarget::beginRenderPass(WGPUCommandEncoder commandEncoder) +{ + assert(commandEncoder); // render pass depth stencil attachment WGPURenderPassDepthStencilAttachment depthStencilAttachment{}; - depthStencilAttachment.view = textureViewStencil; - depthStencilAttachment.depthLoadOp = WGPULoadOp_Clear; - depthStencilAttachment.depthStoreOp = WGPUStoreOp_Store; + depthStencilAttachment.view = texViewStencil; + depthStencilAttachment.depthLoadOp = WGPULoadOp_Load; + depthStencilAttachment.depthStoreOp = WGPUStoreOp_Discard; depthStencilAttachment.depthClearValue = 1.0f; depthStencilAttachment.depthReadOnly = false; depthStencilAttachment.stencilLoadOp = WGPULoadOp_Clear; @@ -157,9 +152,9 @@ void WgRenderTarget::beginRenderPass(WGPUCommandEncoder commandEncoder, WGPUText depthStencilAttachment.stencilReadOnly = false; // render pass color attachment WGPURenderPassColorAttachment colorAttachment{}; - colorAttachment.view = colorAttachement; + colorAttachment.view = texViewColor; colorAttachment.resolveTarget = nullptr; - colorAttachment.loadOp = clear ? WGPULoadOp_Clear : WGPULoadOp_Load; + colorAttachment.loadOp = WGPULoadOp_Clear; colorAttachment.clearValue = { 0, 0, 0, 0 }; colorAttachment.storeOp = WGPUStoreOp_Store; // render pass descriptor @@ -173,134 +168,145 @@ void WgRenderTarget::beginRenderPass(WGPUCommandEncoder commandEncoder, WGPUText renderPassDesc.occlusionQuerySet = nullptr; renderPassDesc.timestampWrites = nullptr; // begin render pass - mRenderPassEncoder = wgpuCommandEncoderBeginRenderPass(commandEncoder, &renderPassDesc); + return wgpuCommandEncoderBeginRenderPass(commandEncoder, &renderPassDesc); } -void WgRenderTarget::endRenderPass() +void WgRenderTarget::endRenderPass(WGPURenderPassEncoder renderPassEncoder) { - assert(mRenderPassEncoder); - wgpuRenderPassEncoderEnd(mRenderPassEncoder); - wgpuRenderPassEncoderRelease(mRenderPassEncoder); - mRenderPassEncoder = nullptr; -} - - -void WgRenderTarget::renderShape(WgRenderDataShape* renderData) -{ - assert(renderData); - assert(mRenderPassEncoder); - // draw shape geometry - wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 0); - for (uint32_t i = 0; i < renderData->meshGroupShapes.meshes.count; i++) { - // draw to stencil (first pass) - mPipelines->fillShape.use(mRenderPassEncoder, mBindGroupCanvasWnd, renderData->bindGroupPaint); - renderData->meshGroupShapes.meshes[i]->draw(mRenderPassEncoder); - // fill shape (second pass) - WgRenderSettings& settings = renderData->renderSettingsShape; - if (settings.fillType == WgRenderSettingsType::Solid) - mPipelines->solid.use(mRenderPassEncoder, mBindGroupCanvasWnd, renderData->bindGroupPaint, settings.bindGroupSolid); - else if (settings.fillType == WgRenderSettingsType::Linear) - mPipelines->linear.use(mRenderPassEncoder, mBindGroupCanvasWnd, renderData->bindGroupPaint, settings.bindGroupLinear); - else if (settings.fillType == WgRenderSettingsType::Radial) - mPipelines->radial.use(mRenderPassEncoder, mBindGroupCanvasWnd, renderData->bindGroupPaint, settings.bindGroupRadial); - renderData->meshBBoxShapes.draw(mRenderPassEncoder); - } -} - - -void WgRenderTarget::renderStroke(WgRenderDataShape* renderData) -{ - assert(renderData); - assert(mRenderPassEncoder); - // draw stroke geometry - if (renderData->meshGroupStrokes.meshes.count > 0) { - // draw strokes to stencil (first pass) - wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 255); - for (uint32_t i = 0; i < renderData->meshGroupStrokes.meshes.count; i++) { - mPipelines->fillStroke.use(mRenderPassEncoder, mBindGroupCanvasWnd, renderData->bindGroupPaint); - renderData->meshGroupStrokes.meshes[i]->draw(mRenderPassEncoder); - } - // fill shape (second pass) - wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 0); - WgRenderSettings& settings = renderData->renderSettingsStroke; - if (settings.fillType == WgRenderSettingsType::Solid) - mPipelines->solid.use(mRenderPassEncoder, mBindGroupCanvasWnd, renderData->bindGroupPaint, settings.bindGroupSolid); - else if (settings.fillType == WgRenderSettingsType::Linear) - mPipelines->linear.use(mRenderPassEncoder, mBindGroupCanvasWnd, renderData->bindGroupPaint, settings.bindGroupLinear); - else if (settings.fillType == WgRenderSettingsType::Radial) - mPipelines->radial.use(mRenderPassEncoder, mBindGroupCanvasWnd, renderData->bindGroupPaint, settings.bindGroupRadial); - renderData->meshBBoxStrokes.draw(mRenderPassEncoder); - } -} - - -void WgRenderTarget::renderPicture(WgRenderDataPicture* renderData) -{ - assert(renderData); - assert(mRenderPassEncoder); - if (renderData->meshData.bufferTexCoord) { - wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 0); - mPipelines->image.use( - mRenderPassEncoder, - mBindGroupCanvasWnd, - renderData->bindGroupPaint, - renderData->bindGroupPicture); - renderData->meshData.drawImage(mRenderPassEncoder); - } -} - - -void WgRenderTarget::blit(WgContext& context, WgRenderTarget* renderTargetSrc, WgBindGroupOpacity* mBindGroupOpacity) -{ - assert(mRenderPassEncoder); - mPipelines->blit.use(mRenderPassEncoder, renderTargetSrc->bindGroupTexSampled, *mBindGroupOpacity); - mMeshDataCanvasWnd.drawImage(mRenderPassEncoder); -} - - -void WgRenderTarget::blitColor(WgContext& context, WgRenderTarget* renderTargetSrc) -{ - assert(mRenderPassEncoder); - mPipelines->blitColor.use(mRenderPassEncoder, renderTargetSrc->bindGroupTexSampled); - mMeshDataCanvasWnd.drawImage(mRenderPassEncoder); -} - - -void WgRenderTarget::compose(WgContext& context, WgRenderTarget* renderTargetSrc, WgRenderTarget* renderTargetMsk, CompositeMethod method) -{ - assert(mRenderPassEncoder); - WgPipelineComposition* pipeline = mPipelines->getCompositionPipeline(method); - assert(pipeline); - pipeline->use(mRenderPassEncoder, renderTargetSrc->bindGroupTexSampled, renderTargetMsk->bindGroupTexSampled); - mMeshDataCanvasWnd.drawImage(mRenderPassEncoder); + assert(renderPassEncoder); + wgpuRenderPassEncoderEnd(renderPassEncoder); + wgpuRenderPassEncoderRelease(renderPassEncoder); } //***************************************************************************** -// render terget pool +// render storage //***************************************************************************** -WgRenderTarget* WgRenderTargetPool::allocate(WgContext& context, uint32_t w, uint32_t h) + void WgRenderStorage::initialize(WgContext& context, uint32_t w, uint32_t h) + { + release(context); + // store target storage size + width = w; + height = h; + workgroupsCountX = (width + WG_COMPUTE_WORKGROUP_SIZE_X - 1) / WG_COMPUTE_WORKGROUP_SIZE_X; // workgroup size x == 8 + workgroupsCountY = (height + WG_COMPUTE_WORKGROUP_SIZE_Y - 1) / WG_COMPUTE_WORKGROUP_SIZE_Y; // workgroup size y == 8 + // create color and stencil textures + texStorage = context.createTexture2d( + WGPUTextureUsage_StorageBinding | WGPUTextureUsage_CopySrc, + WGPUTextureFormat_RGBA8Unorm, + w, h, "The target texture storage color"); + assert(texStorage); + texViewStorage = context.createTextureView2d(texStorage, "The target texture storage view color"); + assert(texViewStorage); + // initialize bind group for blitting + bindGroupTexStorage.initialize(context.device, context.queue, texViewStorage); + mPipelines = context.pipelines; + } + + +void WgRenderStorage::release(WgContext& context) { - WgRenderTarget* renderTarget{}; + bindGroupTexStorage.release(); + context.releaseTextureView(texViewStorage); + context.releaseTexture(texStorage); + workgroupsCountX = 0; + workgroupsCountY = 0; + height = 0; + width = 0; +} + +void WgRenderStorage::clear(WGPUCommandEncoder commandEncoder) +{ + assert(commandEncoder); + WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder); + mPipelines->computeClear.use(computePassEncoder, bindGroupTexStorage); + dispatchWorkgroups(computePassEncoder); + endRenderPass(computePassEncoder); +} + + +void WgRenderStorage::blend(WGPUCommandEncoder commandEncoder, WgRenderTarget* targetSrc, WgBindGroupBlendMethod* blendMethod) +{ + assert(commandEncoder); + assert(targetSrc); + WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder); + mPipelines->computeBlend.use(computePassEncoder, targetSrc->bindGroupTexStorage, bindGroupTexStorage, *blendMethod); + dispatchWorkgroups(computePassEncoder); + endRenderPass(computePassEncoder); +}; + +void WgRenderStorage::blend(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc, WgBindGroupBlendMethod* blendMethod) +{ + assert(commandEncoder); + assert(targetSrc); + WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder); + mPipelines->computeBlend.use(computePassEncoder, targetSrc->bindGroupTexStorage, bindGroupTexStorage, *blendMethod); + dispatchWorkgroups(computePassEncoder); + endRenderPass(computePassEncoder); +}; + + +void WgRenderStorage::compose(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetMsk, WgBindGroupCompositeMethod* composeMethod, WgBindGroupOpacity* opacity) +{ + assert(commandEncoder); + assert(targetMsk); + WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder); + mPipelines->computeCompose.use(computePassEncoder, bindGroupTexStorage, targetMsk->bindGroupTexStorage, *composeMethod, *opacity); + dispatchWorkgroups(computePassEncoder); + endRenderPass(computePassEncoder); +}; + + +void WgRenderStorage::dispatchWorkgroups(WGPUComputePassEncoder computePassEncoder) +{ + assert(computePassEncoder); + wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, workgroupsCountX, workgroupsCountY, 1); +} + + +WGPUComputePassEncoder WgRenderStorage::beginComputePass(WGPUCommandEncoder commandEncoder) +{ + assert(commandEncoder); + WGPUComputePassDescriptor computePassDesc{}; + computePassDesc.nextInChain = nullptr; + computePassDesc.label = "The compute pass composition"; + computePassDesc.timestampWrites = nullptr; + return wgpuCommandEncoderBeginComputePass(commandEncoder, &computePassDesc); +}; + + +void WgRenderStorage::endRenderPass(WGPUComputePassEncoder computePassEncoder) +{ + assert(computePassEncoder); + wgpuComputePassEncoderEnd(computePassEncoder); +} + +//***************************************************************************** +// render storage pool +//***************************************************************************** + +WgRenderStorage* WgRenderStoragePool::allocate(WgContext& context, uint32_t w, uint32_t h) +{ + WgRenderStorage* renderStorage{}; if (mPool.count > 0) { - renderTarget = mPool.last(); + renderStorage = mPool.last(); mPool.pop(); } else { - renderTarget = new WgRenderTarget; - renderTarget->initialize(context, w, h); - mList.push(renderTarget); + renderStorage = new WgRenderStorage; + renderStorage->initialize(context, w, h); + mList.push(renderStorage); } - return renderTarget; + return renderStorage; }; -void WgRenderTargetPool::free(WgContext& context, WgRenderTarget* renderTarget) { - mPool.push(renderTarget); +void WgRenderStoragePool::free(WgContext& context, WgRenderStorage* renderStorage) { + mPool.push(renderStorage); }; -void WgRenderTargetPool::release(WgContext& context) +void WgRenderStoragePool::release(WgContext& context) { for (uint32_t i = 0; i < mList.count; i++) { mList[i]->release(context); diff --git a/src/renderer/wg_engine/tvgWgRenderTarget.h b/src/renderer/wg_engine/tvgWgRenderTarget.h index f81b87d8..a4c85b53 100644 --- a/src/renderer/wg_engine/tvgWgRenderTarget.h +++ b/src/renderer/wg_engine/tvgWgRenderTarget.h @@ -27,48 +27,65 @@ class WgRenderTarget { private: - // command buffer - WGPURenderPassEncoder mRenderPassEncoder{}; - // fill and blit data - WgBindGroupCanvas mBindGroupCanvasWnd; - // composition handles - WgMeshData mMeshDataCanvasWnd; - // gpu buffers - WGPUTexture mTextureColor{}; - WGPUTexture mTextureStencil{}; + // canvas info + WgBindGroupCanvas mBindGroupCanvas; WgPipelines* mPipelines{}; // external handle public: - WGPUSampler sampler{}; - WGPUTextureView textureViewColor{}; - WGPUTextureView textureViewStencil{}; - WgBindGroupTexture bindGroupTex; - WgBindGroupStorageTexture bindGroupStorageTex; - WgBindGroupTextureSampled bindGroupTexSampled; + WGPUTexture texColor{}; + WGPUTexture texStencil{}; + WGPUTextureView texViewColor{}; + WGPUTextureView texViewStencil{}; + WgBindGroupTextureStorage bindGroupTexStorage; public: void initialize(WgContext& context, uint32_t w, uint32_t h); void release(WgContext& context); - void beginRenderPass(WGPUCommandEncoder commandEncoder, WGPUTextureView colorAttachement, bool clear); - void beginRenderPass(WGPUCommandEncoder commandEncoder, bool clear); - void endRenderPass(); + void renderShape(WGPUCommandEncoder commandEncoder, WgRenderDataShape* renderData); + void renderPicture(WGPUCommandEncoder commandEncoder, WgRenderDataPicture* renderData); +private: + void drawShape(WGPURenderPassEncoder renderPassEncoder, WgRenderDataShape* renderData); + void drawStroke(WGPURenderPassEncoder renderPassEncoder, WgRenderDataShape* renderData); - void renderShape(WgRenderDataShape* renderData); - void renderStroke(WgRenderDataShape* renderData); - void renderPicture(WgRenderDataPicture* renderData); - - void blit(WgContext& context, WgRenderTarget* renderTargetSrc, WgBindGroupOpacity* bindGroupOpacity); - void blitColor(WgContext& context, WgRenderTarget* renderTargetSrc); - void compose(WgContext& context, WgRenderTarget* renderTargetSrc, WgRenderTarget* renderTargetMsk, CompositeMethod method); + WGPURenderPassEncoder beginRenderPass(WGPUCommandEncoder commandEncoder); + void endRenderPass(WGPURenderPassEncoder renderPassEncoder); }; -class WgRenderTargetPool { +class WgRenderStorage { private: - Array mList; - Array mPool; + // texture buffers + WgPipelines* mPipelines{}; // external handle public: - WgRenderTarget* allocate(WgContext& context, uint32_t w, uint32_t h); - void free(WgContext& context, WgRenderTarget* renderTarget); + WGPUTexture texStorage{}; + WGPUTextureView texViewStorage{}; + WgBindGroupTextureStorage bindGroupTexStorage; + uint32_t width{}; + uint32_t height{}; + uint32_t workgroupsCountX{}; + uint32_t workgroupsCountY{}; +public: + void initialize(WgContext& context, uint32_t w, uint32_t h); + void release(WgContext& context); + + void clear(WGPUCommandEncoder commandEncoder); + void blend(WGPUCommandEncoder commandEncoder, WgRenderTarget* targetSrc, WgBindGroupBlendMethod* blendMethod); + void blend(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc, WgBindGroupBlendMethod* blendMethod); + void compose(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetMsk, WgBindGroupCompositeMethod* composeMethod, WgBindGroupOpacity* opacity); +private: + void dispatchWorkgroups(WGPUComputePassEncoder computePassEncoder); + + WGPUComputePassEncoder beginComputePass(WGPUCommandEncoder commandEncoder); + void endRenderPass(WGPUComputePassEncoder computePassEncoder); +}; + + +class WgRenderStoragePool { +private: + Array mList; + Array mPool; +public: + WgRenderStorage* allocate(WgContext& context, uint32_t w, uint32_t h); + void free(WgContext& context, WgRenderStorage* renderTarget); void release(WgContext& context); }; diff --git a/src/renderer/wg_engine/tvgWgRenderer.cpp b/src/renderer/wg_engine/tvgWgRenderer.cpp index e6672e3d..0172e9f4 100644 --- a/src/renderer/wg_engine/tvgWgRenderer.cpp +++ b/src/renderer/wg_engine/tvgWgRenderer.cpp @@ -43,18 +43,22 @@ void WgRenderer::initialize() { mContext.initialize(); mPipelines.initialize(mContext); - mBindGroupOpacityPool.initialize(mContext); + mOpacityPool.initialize(mContext); + mBlendMethodPool.initialize(mContext); + mCompositeMethodPool.initialize(mContext); } void WgRenderer::release() { mCompositorStack.clear(); - mRenderTargetStack.clear(); - mBindGroupOpacityPool.release(mContext); - mRenderTargetPool.release(mContext); - mRenderTargetRoot.release(mContext); - mRenderTargetWnd.release(mContext); + mRenderStorageStack.clear(); + mRenderStoragePool.release(mContext); + mCompositeMethodPool.release(mContext); + mBlendMethodPool.release(mContext); + mOpacityPool.release(mContext); + mRenderStorageRoot.release(mContext); + mRenderTarget.release(mContext); wgpuSurfaceUnconfigure(mSurface); wgpuSurfaceRelease(mSurface); mPipelines.release(); @@ -108,7 +112,7 @@ RenderData WgRenderer::prepare(Surface* surface, const RenderMesh* mesh, RenderD WgShaderTypeBlendSettings blendSettings(surface->cs, opacity); renderDataPicture->bindGroupPaint.initialize(mContext.device, mContext.queue, modelMat, blendSettings); } - + // update image data if (flags & (RenderUpdateFlag::Path | RenderUpdateFlag::Image)) { WgGeometryData geometryData; @@ -119,7 +123,7 @@ RenderData WgRenderer::prepare(Surface* surface, const RenderMesh* mesh, RenderD renderDataPicture->imageData.update(mContext, surface); renderDataPicture->bindGroupPicture.initialize( mContext.device, mContext.queue, - renderDataPicture->imageData.sampler, + mContext.samplerLinear, renderDataPicture->imageData.textureView); } @@ -129,37 +133,41 @@ RenderData WgRenderer::prepare(Surface* surface, const RenderMesh* mesh, RenderD bool WgRenderer::preRender() { - // command encoder descriptor WGPUCommandEncoderDescriptor commandEncoderDesc{}; commandEncoderDesc.nextInChain = nullptr; commandEncoderDesc.label = "The command encoder"; mCommandEncoder = wgpuDeviceCreateCommandEncoder(mContext.device, &commandEncoderDesc); - // render datas - mRenderTargetStack.push(&mRenderTargetRoot); - mRenderTargetRoot.beginRenderPass(mCommandEncoder, true); + mRenderStorageRoot.clear(mCommandEncoder); + mRenderStorageStack.push(&mRenderStorageRoot); return true; } bool WgRenderer::renderShape(RenderData data) { - mRenderTargetStack.last()->renderShape((WgRenderDataShape *)data); - mRenderTargetStack.last()->renderStroke((WgRenderDataShape *)data); + // render shape to render target + mRenderTarget.renderShape(mCommandEncoder, (WgRenderDataShape *)data); + // blend shape with current render storage + WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod); + mRenderStorageStack.last()->blend(mCommandEncoder, &mRenderTarget, blendMethod); return true; } bool WgRenderer::renderImage(RenderData data) { - mRenderTargetStack.last()->renderPicture((WgRenderDataPicture *)data); + // render image to render target + mRenderTarget.renderPicture(mCommandEncoder, (WgRenderDataPicture *)data); + // blend image with current render storage + WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod); + mRenderStorageStack.last()->blend(mCommandEncoder, &mRenderTarget, blendMethod); return true; } bool WgRenderer::postRender() { - mRenderTargetRoot.endRenderPass(); - mRenderTargetStack.pop(); + mRenderStorageStack.pop(); mContext.executeCommandEncoder(mCommandEncoder); wgpuCommandEncoderRelease(mCommandEncoder); return true; @@ -193,6 +201,7 @@ bool WgRenderer::viewport(TVG_UNUSED const RenderRegion& vp) bool WgRenderer::blend(TVG_UNUSED BlendMethod method) { + mBlendMethod = method; return false; } @@ -213,17 +222,25 @@ bool WgRenderer::sync() { WGPUSurfaceTexture backBuffer{}; wgpuSurfaceGetCurrentTexture(mSurface, &backBuffer); - WGPUTextureView backBufferView = mContext.createTextureView2d(backBuffer.texture, "Surface texture view"); + WGPUCommandEncoderDescriptor commandEncoderDesc{}; commandEncoderDesc.nextInChain = nullptr; commandEncoderDesc.label = "The command encoder"; WGPUCommandEncoder commandEncoder = wgpuDeviceCreateCommandEncoder(mContext.device, &commandEncoderDesc); - mRenderTargetWnd.beginRenderPass(commandEncoder, backBufferView, true); - mRenderTargetWnd.blitColor(mContext, &mRenderTargetRoot); - mRenderTargetWnd.endRenderPass(); + + WGPUImageCopyTexture source{}; + source.texture = mRenderStorageRoot.texStorage; + WGPUImageCopyTexture dest{}; + dest.texture = backBuffer.texture; + WGPUExtent3D copySize{}; + copySize.width = mTargetSurface.w; + copySize.height = mTargetSurface.h; + copySize.depthOrArrayLayers = 1; + wgpuCommandEncoderCopyTextureToTexture(commandEncoder, &source, &dest, ©Size); + mContext.executeCommandEncoder(commandEncoder); wgpuCommandEncoderRelease(commandEncoder); - wgpuTextureViewRelease(backBufferView); + wgpuSurfacePresent(mSurface); return true; } @@ -236,7 +253,7 @@ bool WgRenderer::target(uint32_t* buffer, uint32_t stride, uint32_t w, uint32_t mTargetSurface.w = w; mTargetSurface.h = h; - mRenderTargetRoot.initialize(mContext, w, h); + mRenderTarget.initialize(mContext, w, h); return true; } @@ -266,7 +283,7 @@ bool WgRenderer::target(void* window, uint32_t w, uint32_t h) surfaceConfiguration.nextInChain = nullptr; surfaceConfiguration.device = mContext.device; surfaceConfiguration.format = WGPUTextureFormat_RGBA8Unorm; - surfaceConfiguration.usage = WGPUTextureUsage_RenderAttachment; + surfaceConfiguration.usage = WGPUTextureUsage_RenderAttachment | WGPUTextureUsage_CopyDst; surfaceConfiguration.viewFormatCount = 0; surfaceConfiguration.viewFormats = nullptr; surfaceConfiguration.alphaMode = WGPUCompositeAlphaMode_Auto; @@ -275,8 +292,8 @@ bool WgRenderer::target(void* window, uint32_t w, uint32_t h) surfaceConfiguration.presentMode = WGPUPresentMode_Mailbox; wgpuSurfaceConfigure(mSurface, &surfaceConfiguration); - mRenderTargetWnd.initialize(mContext, w, h); - mRenderTargetRoot.initialize(mContext, w, h); + mRenderTarget.initialize(mContext, w, h); + mRenderStorageRoot.initialize(mContext, w, h); return true; } @@ -291,17 +308,14 @@ Compositor* WgRenderer::target(TVG_UNUSED const RenderRegion& region, TVG_UNUSED bool WgRenderer::beginComposite(TVG_UNUSED Compositor* cmp, TVG_UNUSED CompositeMethod method, TVG_UNUSED uint8_t opacity) { - // save current composition settings + // save current composition settings cmp->method = method; cmp->opacity = opacity; - // end current render target - mRenderTargetStack.last()->endRenderPass(); - - // create new render target and begin new render pass - WgRenderTarget* renderTarget = mRenderTargetPool.allocate(mContext, mTargetSurface.w, mTargetSurface.h); - renderTarget->beginRenderPass(mCommandEncoder, true); - mRenderTargetStack.push(renderTarget); + // allocate new render storage and push it to top of render tree + WgRenderStorage* renderStorage = mRenderStoragePool.allocate(mContext, mTargetSurface.w, mTargetSurface.h); + renderStorage->clear(mCommandEncoder); + mRenderStorageStack.push(renderStorage); return true; } @@ -310,41 +324,35 @@ bool WgRenderer::beginComposite(TVG_UNUSED Compositor* cmp, TVG_UNUSED Composite bool WgRenderer::endComposite(TVG_UNUSED Compositor* cmp) { if (cmp->method == CompositeMethod::None) { - // end current render pass - mRenderTargetStack.last()->endRenderPass(); - // get two last render targets - WgRenderTarget* renderTargetSrc = mRenderTargetStack.last(); - mRenderTargetStack.pop(); + WgRenderStorage* renderStorageSrc = mRenderStorageStack.last(); + mRenderStorageStack.pop(); - // apply current render target - WgRenderTarget* renderTarget = mRenderTargetStack.last(); - renderTarget->beginRenderPass(mCommandEncoder, false); - - // blit scene to current render tartget - WgBindGroupOpacity* mBindGroupOpacity = mBindGroupOpacityPool.allocate(mContext, cmp->opacity); - renderTarget->blit(mContext, renderTargetSrc, mBindGroupOpacity); + // blent scene to current render storage + WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod); + mRenderStorageStack.last()->blend(mCommandEncoder, renderStorageSrc, blendMethod); // back render targets to the pool - mRenderTargetPool.free(mContext, renderTargetSrc); + mRenderStoragePool.free(mContext, renderStorageSrc); } else { - // end current render pass - mRenderTargetStack.last()->endRenderPass(); - // get two last render targets - WgRenderTarget* renderTargetSrc = mRenderTargetStack.last(); - mRenderTargetStack.pop(); - WgRenderTarget* renderTargetMsk = mRenderTargetStack.last(); - mRenderTargetStack.pop(); + WgRenderStorage* renderStorageSrc = mRenderStorageStack.last(); + mRenderStorageStack.pop(); + WgRenderStorage* renderStorageMsk = mRenderStorageStack.last(); + mRenderStorageStack.pop(); - // apply current render target - WgRenderTarget* renderTarget = mRenderTargetStack.last(); - renderTarget->beginRenderPass(mCommandEncoder, false); - renderTarget->compose(mContext, renderTargetSrc, renderTargetMsk, cmp->method); + // compose shape and mask + WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, cmp->opacity); + WgBindGroupCompositeMethod* composeMethod = mCompositeMethodPool.allocate(mContext, cmp->method); + renderStorageSrc->compose(mCommandEncoder, renderStorageMsk, composeMethod, opacity); + + // blent scene to current render storage + WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod); + mRenderStorageStack.last()->blend(mCommandEncoder, renderStorageSrc, blendMethod); // back render targets to the pool - mRenderTargetPool.free(mContext, renderTargetSrc); - mRenderTargetPool.free(mContext, renderTargetMsk); + mRenderStoragePool.free(mContext, renderStorageSrc); + mRenderStoragePool.free(mContext, renderStorageMsk); } // delete current compositor diff --git a/src/renderer/wg_engine/tvgWgRenderer.h b/src/renderer/wg_engine/tvgWgRenderer.h index 3def37f5..ab76a83e 100644 --- a/src/renderer/wg_engine/tvgWgRenderer.h +++ b/src/renderer/wg_engine/tvgWgRenderer.h @@ -62,19 +62,24 @@ public: private: // render handles WGPUCommandEncoder mCommandEncoder{}; - Array mCompositorStack; - Array mRenderTargetStack; + WgRenderTarget mRenderTarget; + WgRenderStorage mRenderStorageRoot; + WgRenderStoragePool mRenderStoragePool; + WgBindGroupOpacityPool mOpacityPool; + WgBindGroupBlendMethodPool mBlendMethodPool; + WgBindGroupCompositeMethodPool mCompositeMethodPool; - // render object pools - WgRenderTargetPool mRenderTargetPool; - WgBindGroupOpacityPool mBindGroupOpacityPool; + // render tree stacks + Array mCompositorStack; + Array mRenderStorageStack; + + // native window handles + WGPUSurface mSurface{}; private: WgContext mContext; WgPipelines mPipelines; - WgRenderTarget mRenderTargetRoot; - WgRenderTarget mRenderTargetWnd; - WGPUSurface mSurface{}; Surface mTargetSurface; + BlendMethod mBlendMethod{}; }; #endif /* _TVG_WG_RENDERER_H_ */ diff --git a/src/renderer/wg_engine/tvgWgShaderSrc.cpp b/src/renderer/wg_engine/tvgWgShaderSrc.cpp index 49efb59c..c3e423a6 100644 --- a/src/renderer/wg_engine/tvgWgShaderSrc.cpp +++ b/src/renderer/wg_engine/tvgWgShaderSrc.cpp @@ -712,18 +712,92 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f { // cShaderSource_PipelineComputeBlend //************************************************************************ -// pipeline shader modules blend (simple example) -const char* cShaderSource_PipelineComputeBlend = R"( -@group(0) @binding(0) var imageSrc : texture_storage_2d; -@group(1) @binding(0) var imageDst : texture_storage_2d; +// pipeline shader modules clear +const char* cShaderSource_PipelineComputeClear = R"( +@group(0) @binding(0) var imageDst : texture_storage_2d; @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; }; - - // var src = textureLoad(imageSrc, id.xy); - textureStore(imageDst, id.xy, vec4(0.5, 1.0, 0.1, 1.0)); + textureStore(imageDst, id.xy, vec4f(0.0, 0.0, 0.0, 0.0)); +} +)"; + + +// pipeline shader modules blend +const char* cShaderSource_PipelineComputeBlend = R"( +@group(0) @binding(0) var imageSrc : texture_storage_2d; +@group(1) @binding(0) var imageDst : texture_storage_2d; +@group(2) @binding(0) var blendMethod : u32; + +@compute @workgroup_size(8, 8) +fn cs_main( @builtin(global_invocation_id) id: vec3u) { + let texSize = textureDimensions(imageSrc); + if ((id.x >= texSize.x) || (id.y >= texSize.y)) { return; }; + + let srcColor = textureLoad(imageSrc, id.xy); + if (srcColor.a == 0.0) { return; }; + let dstColor = textureLoad(imageDst, id.xy); + + let Sa: f32 = srcColor.a; + let Da: f32 = dstColor.a; + let S: vec3f = srcColor.xyz; + let D: vec3f = dstColor.xyz; + let One: vec3f = vec3(1.0); + + var color: vec3f = vec3f(0.0, 0.0, 0.0); + switch blendMethod { + /* Normal */ case 0u: { color = (Sa * S) + (1.0 - Sa) * D; } + /* Add */ case 1u: { color = (S + D); } + /* Screen */ case 2u: { color = (S + D) - (S * D); } + /* Multiply */ case 3u: { color = (S * D); } + /* Overlay */ case 4u: { color = S * D; } + /* Difference */ case 5u: { color = abs(S - D); } + /* Exclusion */ case 6u: { color = S + D - (2 * S * D); } + /* SrcOver */ case 7u: { color = S; } + /* Darken */ case 8u: { color = min(S, D); } + /* Lighten */ case 9u: { color = max(S, D); } + /* ColorDodge */ case 10u: { color = D / (One - S); } + /* ColorBurn */ case 11u: { color = One - (One - D) / S; } + /* HardLight */ case 12u: { color = (Sa * Da) - 2.0 * (Da - S) * (Sa - D); } + /* SoftLight */ case 13u: { color = (One - 2 * S) * (D * D) + (2 * S * D); } + default: { color = (Sa * S) + (1.0 - Sa) * D; } + } + + textureStore(imageDst, id.xy, vec4f(color, Sa)); +} +)"; + +// pipeline shader modules compose +const char* cShaderSource_PipelineComputeCompose = R"( +@group(0) @binding(0) var imageSrc : texture_storage_2d; +@group(1) @binding(0) var imageMsk : texture_storage_2d; +@group(2) @binding(0) var composeMethod : u32; +@group(3) @binding(0) var opacity : f32; + +@compute @workgroup_size(8, 8) +fn cs_main( @builtin(global_invocation_id) id: vec3u) { + let texSize = textureDimensions(imageSrc); + if ((id.x >= texSize.x) || (id.y >= texSize.y)) { return; }; + + let colorSrc = textureLoad(imageSrc, id.xy); + let colorMsk = textureLoad(imageMsk, id.xy); + + var color: vec3f = colorSrc.xyz; + var alpha: f32 = colorMsk.a; + switch composeMethod { + /* None */ case 0u: { color = colorSrc.xyz; } + /* ClipPath */ case 1u: { if (colorMsk.a == 0) { alpha = 0.0; }; } + /* AlphaMask */ case 2u: { color = mix(colorMsk.xyz, colorSrc.xyz, colorSrc.a * colorMsk.b); } + /* InvAlphaMask */ case 3u: { color = mix(colorSrc.xyz, colorMsk.xyz, colorSrc.a * colorMsk.b); alpha = 1.0 - colorMsk.b; } + /* LumaMask */ case 4u: { color = colorSrc.xyz * (0.299 * colorMsk.r + 0.587 * colorMsk.g + 0.114 * colorMsk.b); } + /* InvLumaMask */ case 5u: { color = colorSrc.xyz * (1.0 - (0.299 * colorMsk.r + 0.587 * colorMsk.g + 0.114 * colorMsk.b)); alpha = 1.0 - colorMsk.b; } + /* AddMask */ case 6u: { color = colorSrc.xyz * colorSrc.a + colorMsk.xyz * (1.0 - colorSrc.a); } + /* SubtractMask */ case 7u: { color = colorSrc.xyz * colorSrc.a - colorMsk.xyz * (1.0 - colorSrc.a); } + /* IntersectMask */ case 8u: { color = colorSrc.xyz * min(colorSrc.a, colorMsk.a); } + /* DifferenceMask */ case 9u: { color = abs(colorSrc.xyz - colorMsk.xyz * (1.0 - colorMsk.a)); } + default: { color = colorSrc.xyz; } + } + + textureStore(imageSrc, id.xy, vec4f(color, alpha * opacity)); } - )"; diff --git a/src/renderer/wg_engine/tvgWgShaderSrc.h b/src/renderer/wg_engine/tvgWgShaderSrc.h index 448b85d2..4fbfa646 100644 --- a/src/renderer/wg_engine/tvgWgShaderSrc.h +++ b/src/renderer/wg_engine/tvgWgShaderSrc.h @@ -54,7 +54,9 @@ extern const char* cShaderSource_PipelineCompDifferenceMask; // compute shader modules //***************************************************************************** -// pipeline shader modules blend +// pipeline shader modules clear, compose and blend +extern const char* cShaderSource_PipelineComputeClear; extern const char* cShaderSource_PipelineComputeBlend; +extern const char* cShaderSource_PipelineComputeCompose; #endif // _TVG_WG_SHADER_SRC_H_