From 84d3ef7184d4a5388c94d936dd144e37b380e20a Mon Sep 17 00:00:00 2001 From: Sergii Liebodkin Date: Tue, 23 Jan 2024 11:08:26 +0200 Subject: [PATCH] wg_engine: introduced compute pipeline entities introduces posibility to create compute pipelines does not affect functionality --- src/renderer/wg_engine/tvgWgBindGroups.cpp | 142 +++++++++++++----- src/renderer/wg_engine/tvgWgBindGroups.h | 47 ++++-- src/renderer/wg_engine/tvgWgCommon.cpp | 145 +++++++++++++++++-- src/renderer/wg_engine/tvgWgCommon.h | 22 ++- src/renderer/wg_engine/tvgWgPipelines.cpp | 41 +++++- src/renderer/wg_engine/tvgWgPipelines.h | 45 +++++- src/renderer/wg_engine/tvgWgRenderTarget.cpp | 19 ++- src/renderer/wg_engine/tvgWgRenderTarget.h | 4 +- src/renderer/wg_engine/tvgWgRenderer.cpp | 40 ++--- src/renderer/wg_engine/tvgWgRenderer.h | 3 +- src/renderer/wg_engine/tvgWgShaderSrc.cpp | 20 +++ src/renderer/wg_engine/tvgWgShaderSrc.h | 25 ++-- 12 files changed, 432 insertions(+), 121 deletions(-) diff --git a/src/renderer/wg_engine/tvgWgBindGroups.cpp b/src/renderer/wg_engine/tvgWgBindGroups.cpp index 625bd5ee..3e194b16 100644 --- a/src/renderer/wg_engine/tvgWgBindGroups.cpp +++ b/src/renderer/wg_engine/tvgWgBindGroups.cpp @@ -22,15 +22,20 @@ #include "tvgWgBindGroups.h" - +// canvas information group WGPUBindGroupLayout WgBindGroupCanvas::layout = nullptr; +// paint object information group WGPUBindGroupLayout WgBindGroupPaint::layout = nullptr; +// fill properties information groups WGPUBindGroupLayout WgBindGroupSolidColor::layout = nullptr; WGPUBindGroupLayout WgBindGroupLinearGradient::layout = nullptr; WGPUBindGroupLayout WgBindGroupRadialGradient::layout = nullptr; WGPUBindGroupLayout WgBindGroupPicture::layout = nullptr; +// composition and blending properties gropus WGPUBindGroupLayout WgBindGroupOpacity::layout = nullptr; -WGPUBindGroupLayout WgBindGroupBlit::layout = nullptr; +WGPUBindGroupLayout WgBindGroupTexture::layout = nullptr; +WGPUBindGroupLayout WgBindGroupStorageTexture::layout = nullptr; +WGPUBindGroupLayout WgBindGroupTextureSampled::layout = nullptr; WGPUBindGroupLayout WgBindGroupCanvas::getLayout(WGPUDevice device) @@ -227,7 +232,7 @@ WGPUBindGroupLayout WgBindGroupPicture::getLayout(WGPUDevice device) if (layout) return layout; const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { makeBindGroupLayoutEntrySampler(0), - makeBindGroupLayoutEntryTextureView(1) + makeBindGroupLayoutEntryTexture(1) }; layout = createBindGroupLayout(device, bindGroupLayoutEntries, 2); assert(layout); @@ -303,12 +308,82 @@ void WgBindGroupOpacity::release() } -WGPUBindGroupLayout WgBindGroupBlit::getLayout(WGPUDevice device) +WGPUBindGroupLayout WgBindGroupTexture::getLayout(WGPUDevice device) +{ + if (layout) return layout; + const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { + makeBindGroupLayoutEntryTexture(0) + }; + layout = createBindGroupLayout(device, bindGroupLayoutEntries, 1); + assert(layout); + return layout; +} + + +void WgBindGroupTexture::releaseLayout() +{ + releaseBindGroupLayout(layout); +} + + +void WgBindGroupTexture::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexture) +{ + release(); + const WGPUBindGroupEntry bindGroupEntries[] { + makeBindGroupEntryTextureView(0, uTexture) + }; + mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 1); + assert(mBindGroup); +} + + +void WgBindGroupTexture::release() +{ + releaseBindGroup(mBindGroup); +} + + +WGPUBindGroupLayout WgBindGroupStorageTexture::getLayout(WGPUDevice device) +{ + if (layout) return layout; + const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { + makeBindGroupLayoutEntryStorageTexture(0) + }; + layout = createBindGroupLayout(device, bindGroupLayoutEntries, 1); + assert(layout); + return layout; +} + + +void WgBindGroupStorageTexture::releaseLayout() +{ + releaseBindGroupLayout(layout); +} + + +void WgBindGroupStorageTexture::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexture) +{ + release(); + const WGPUBindGroupEntry bindGroupEntries[] { + makeBindGroupEntryTextureView(0, uTexture) + }; + mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 1); + assert(mBindGroup); +} + + +void WgBindGroupStorageTexture::release() +{ + releaseBindGroup(mBindGroup); +} + + +WGPUBindGroupLayout WgBindGroupTextureSampled::getLayout(WGPUDevice device) { if (layout) return layout; const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { makeBindGroupLayoutEntrySampler(0), - makeBindGroupLayoutEntryTextureView(1) + makeBindGroupLayoutEntryTexture(1) }; layout = createBindGroupLayout(device, bindGroupLayoutEntries, 2); assert(layout); @@ -316,13 +391,13 @@ WGPUBindGroupLayout WgBindGroupBlit::getLayout(WGPUDevice device) } -void WgBindGroupBlit::releaseLayout() +void WgBindGroupTextureSampled::releaseLayout() { releaseBindGroupLayout(layout); } -void WgBindGroupBlit::initialize(WGPUDevice device, WGPUQueue queue, WGPUSampler uSampler, WGPUTextureView uTexture) +void WgBindGroupTextureSampled::initialize(WGPUDevice device, WGPUQueue queue, WGPUSampler uSampler, WGPUTextureView uTexture) { release(); const WGPUBindGroupEntry bindGroupEntries[] { @@ -334,49 +409,40 @@ void WgBindGroupBlit::initialize(WGPUDevice device, WGPUQueue queue, WGPUSampler } -void WgBindGroupBlit::release() +void WgBindGroupTextureSampled::release() { releaseBindGroup(mBindGroup); } //************************************************************************ -// bind groups pools +// bind group pools //************************************************************************ -WgBindGroupOpacity* WgBindGroupOpacityPool::allocate(WgContext& context, uint32_t opacity) +void WgBindGroupOpacityPool::initialize(WgContext& context) { - WgBindGroupOpacity* bindGroup{}; - if (mPool.count == 0) { - bindGroup = new WgBindGroupOpacity; - bindGroup->initialize(context.device, context.queue, opacity); - mUsed.push(bindGroup); - mList.push(bindGroup); - } else { - bindGroup = mPool.last(); - bindGroup->update(context.device, context.queue, opacity); - mUsed.push(bindGroup); - mPool.pop(); + memset(mPool, 0x00, sizeof(mPool)); +} + + +void WgBindGroupOpacityPool::release(WgContext& context) +{ + for (uint32_t i = 0; i < 256; i++) { + if (mPool[i]) { + mPool[i]->release(); + delete mPool[i]; + mPool[i] = nullptr; + } } - return bindGroup; } -void WgBindGroupOpacityPool::reset() +WgBindGroupOpacity* WgBindGroupOpacityPool::allocate(WgContext& context, uint8_t opacity) { - for (uint32_t i = 0; i < mUsed.count; i++) - mPool.push(mUsed[i]); - mUsed.clear(); -} - - -void WgBindGroupOpacityPool::release() -{ - for (uint32_t i = 0; i < mList.count; i++) { - mList[i]->release(); - delete mList[i]; + WgBindGroupOpacity* bindGroupOpacity = mPool[opacity]; + if (!bindGroupOpacity) { + bindGroupOpacity = new WgBindGroupOpacity; + bindGroupOpacity->initialize(context.device, context.queue, opacity); + mPool[opacity] = bindGroupOpacity; } - mList.clear(); - mUsed.clear(); - mPool.clear(); + return bindGroupOpacity; } - diff --git a/src/renderer/wg_engine/tvgWgBindGroups.h b/src/renderer/wg_engine/tvgWgBindGroups.h index b8bc0d16..42c06c00 100644 --- a/src/renderer/wg_engine/tvgWgBindGroups.h +++ b/src/renderer/wg_engine/tvgWgBindGroups.h @@ -39,7 +39,6 @@ struct WgBindGroupCanvas : public WgBindGroup void release(); }; - // @group(1) struct WgBindGroupPaint : public WgBindGroup { @@ -56,7 +55,7 @@ struct WgBindGroupPaint : public WgBindGroup }; // @group(2) -struct WgBindGroupSolidColor : public WgBindGroup +struct WgBindGroupSolidColor : public WgBindGroup { static WGPUBindGroupLayout layout; static WGPUBindGroupLayout getLayout(WGPUDevice device); @@ -69,7 +68,7 @@ struct WgBindGroupSolidColor : public WgBindGroup }; // @group(2) -struct WgBindGroupLinearGradient : public WgBindGroup +struct WgBindGroupLinearGradient : public WgBindGroup { static WGPUBindGroupLayout layout; static WGPUBindGroupLayout getLayout(WGPUDevice device); @@ -82,7 +81,7 @@ struct WgBindGroupLinearGradient : public WgBindGroup }; // @group(2) -struct WgBindGroupRadialGradient : public WgBindGroup +struct WgBindGroupRadialGradient : public WgBindGroup { static WGPUBindGroupLayout layout; static WGPUBindGroupLayout getLayout(WGPUDevice device); @@ -95,7 +94,7 @@ struct WgBindGroupRadialGradient : public WgBindGroup }; // @group(2) -struct WgBindGroupPicture : public WgBindGroup +struct WgBindGroupPicture : public WgBindGroup { static WGPUBindGroupLayout layout; static WGPUBindGroupLayout getLayout(WGPUDevice device); @@ -108,7 +107,7 @@ struct WgBindGroupPicture : public WgBindGroup }; // @group(1 or 2) -struct WgBindGroupOpacity : public WgBindGroup +struct WgBindGroupOpacity : public WgBindGroup { static WGPUBindGroupLayout layout; static WGPUBindGroupLayout getLayout(WGPUDevice device); @@ -121,7 +120,31 @@ struct WgBindGroupOpacity : public WgBindGroup }; // @group(0 or 1) -struct WgBindGroupBlit : public WgBindGroup +struct WgBindGroupTexture : public WgBindGroup +{ + static WGPUBindGroupLayout layout; + static WGPUBindGroupLayout getLayout(WGPUDevice device); + static void releaseLayout(); + + void initialize(WGPUDevice device, WGPUQueue queue, + WGPUTextureView uTexture); + void release(); +}; + +// @group(0 or 1) +struct WgBindGroupStorageTexture : public WgBindGroup +{ + static WGPUBindGroupLayout layout; + static WGPUBindGroupLayout getLayout(WGPUDevice device); + static void releaseLayout(); + + void initialize(WGPUDevice device, WGPUQueue queue, + WGPUTextureView uTexture); + void release(); +}; + +// @group(0 or 1) +struct WgBindGroupTextureSampled : public WgBindGroup { static WGPUBindGroupLayout layout; static WGPUBindGroupLayout getLayout(WGPUDevice device); @@ -140,13 +163,11 @@ struct WgBindGroupBlit : public WgBindGroup class WgBindGroupOpacityPool { private: - Array mList; - Array mPool; - Array mUsed; + WgBindGroupOpacity* mPool[256]; public: - WgBindGroupOpacity* allocate(WgContext& context, uint32_t opacity); - void reset(); - void release(); + void initialize(WgContext& context); + void release(WgContext& context); + WgBindGroupOpacity* allocate(WgContext& context, uint8_t opacity); }; #endif // _TVG_WG_BIND_GROUPS_H_ diff --git a/src/renderer/wg_engine/tvgWgCommon.cpp b/src/renderer/wg_engine/tvgWgCommon.cpp index 4f00f6fe..e0d6e8d2 100644 --- a/src/renderer/wg_engine/tvgWgCommon.cpp +++ b/src/renderer/wg_engine/tvgWgCommon.cpp @@ -60,7 +60,7 @@ void WgContext::initialize() WGPUDeviceDescriptor deviceDesc{}; deviceDesc.nextInChain = nullptr; deviceDesc.label = "The device"; - deviceDesc.requiredFeaturesCount = featuresCount; + deviceDesc.requiredFeatureCount = featuresCount; deviceDesc.requiredFeatures = featureNames; deviceDesc.requiredLimits = nullptr; deviceDesc.defaultQueue.nextInChain = nullptr; @@ -115,6 +115,54 @@ void WgContext::executeCommandEncoder(WGPUCommandEncoder commandEncoder) } +WGPUTexture WgContext::createTexture2d(WGPUTextureUsage usage, WGPUTextureFormat format, uint32_t width, uint32_t height, char const * label) { + WGPUTextureDescriptor textureDesc{}; + textureDesc.nextInChain = nullptr; + textureDesc.label = label; + textureDesc.usage = usage; + textureDesc.dimension = WGPUTextureDimension_2D; + textureDesc.size = { width, height, 1 }; + textureDesc.format = format; + textureDesc.mipLevelCount = 1; + textureDesc.sampleCount = 1; + textureDesc.viewFormatCount = 0; + textureDesc.viewFormats = nullptr; + return wgpuDeviceCreateTexture(device, &textureDesc); +} + + +WGPUTextureView WgContext::createTextureView2d(WGPUTexture texture, WGPU_NULLABLE char const * label) +{ + WGPUTextureViewDescriptor textureViewDescColor{}; + textureViewDescColor.nextInChain = nullptr; + textureViewDescColor.label = label; + textureViewDescColor.format = wgpuTextureGetFormat(texture); + textureViewDescColor.dimension = WGPUTextureViewDimension_2D; + textureViewDescColor.baseMipLevel = 0; + textureViewDescColor.mipLevelCount = 1; + textureViewDescColor.baseArrayLayer = 0; + textureViewDescColor.arrayLayerCount = 1; + textureViewDescColor.aspect = WGPUTextureAspect_All; + return wgpuTextureCreateView(texture, &textureViewDescColor); +}; + + +void WgContext::releaseTexture(WGPUTexture& texture) { + if (texture) { + wgpuTextureDestroy(texture); + wgpuTextureRelease(texture); + texture = nullptr; + } + +} + + +void WgContext::releaseTextureView(WGPUTextureView& textureView) { + if (textureView) wgpuTextureViewRelease(textureView); + textureView = nullptr; +} + + //***************************************************************************** // bind group //***************************************************************************** @@ -125,6 +173,12 @@ void WgBindGroup::set(WGPURenderPassEncoder encoder, uint32_t groupIndex) } +void WgBindGroup::set(WGPUComputePassEncoder encoder, uint32_t groupIndex) +{ + wgpuComputePassEncoderSetBindGroup(encoder, groupIndex, mBindGroup, 0, nullptr); +} + + WGPUBindGroupEntry WgBindGroup::makeBindGroupEntryBuffer(uint32_t binding, WGPUBuffer buffer) { WGPUBindGroupEntry bindGroupEntry{}; @@ -193,12 +247,12 @@ WGPUBindGroupLayoutEntry WgBindGroup::makeBindGroupLayoutEntrySampler(uint32_t b } -WGPUBindGroupLayoutEntry WgBindGroup::makeBindGroupLayoutEntryTextureView(uint32_t binding) +WGPUBindGroupLayoutEntry WgBindGroup::makeBindGroupLayoutEntryTexture(uint32_t binding) { WGPUBindGroupLayoutEntry bindGroupLayoutEntry{}; bindGroupLayoutEntry.nextInChain = nullptr; bindGroupLayoutEntry.binding = binding; - bindGroupLayoutEntry.visibility = WGPUShaderStage_Fragment; + bindGroupLayoutEntry.visibility = WGPUShaderStage_Fragment | WGPUShaderStage_Compute; bindGroupLayoutEntry.texture.nextInChain = nullptr; bindGroupLayoutEntry.texture.sampleType = WGPUTextureSampleType_Float; bindGroupLayoutEntry.texture.viewDimension = WGPUTextureViewDimension_2D; @@ -207,6 +261,20 @@ WGPUBindGroupLayoutEntry WgBindGroup::makeBindGroupLayoutEntryTextureView(uint32 } +WGPUBindGroupLayoutEntry WgBindGroup::makeBindGroupLayoutEntryStorageTexture(uint32_t binding) +{ + 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.format = WGPUTextureFormat_RGBA8Unorm; + bindGroupLayoutEntry.storageTexture.viewDimension = WGPUTextureViewDimension_2D; + return bindGroupLayoutEntry; +} + + WGPUBuffer WgBindGroup::createBuffer(WGPUDevice device, WGPUQueue queue, const void *data, size_t size) { WGPUBufferDescriptor bufferDescriptor{}; @@ -226,7 +294,7 @@ WGPUBindGroup WgBindGroup::createBindGroup(WGPUDevice device, WGPUBindGroupLayou { WGPUBindGroupDescriptor bindGroupDesc{}; bindGroupDesc.nextInChain = nullptr; - bindGroupDesc.label = "The binding group sampler"; + bindGroupDesc.label = "The binding group"; bindGroupDesc.layout = layout; bindGroupDesc.entryCount = count; bindGroupDesc.entries = bindGroupEntries; @@ -283,7 +351,7 @@ WGPUPipelineLayout WgPipeline::createPipelineLayout(WGPUDevice device, const WGP { WGPUPipelineLayoutDescriptor pipelineLayoutDesc{}; pipelineLayoutDesc.nextInChain = nullptr; - pipelineLayoutDesc.label = "The Pipeline layout"; + pipelineLayoutDesc.label = "The pipeline layout"; pipelineLayoutDesc.bindGroupLayoutCount = count; pipelineLayoutDesc.bindGroupLayouts = bindGroupLayouts; return wgpuDeviceCreatePipelineLayout(device, &pipelineLayoutDesc); @@ -322,11 +390,11 @@ void WgPipeline::destroyShaderModule(WGPUShaderModule& shaderModule) // render pipeline //***************************************************************************** -void WgRenderPipeline::allocate(WGPUDevice device, - WGPUVertexBufferLayout vertexBufferLayouts[], uint32_t attribsCount, - WGPUBindGroupLayout bindGroupLayouts[], uint32_t bindGroupsCount, - WGPUCompareFunction stencilCompareFunction, WGPUStencilOperation stencilOperation, - const char* shaderSource, const char* shaderLabel, const char* pipelineLabel) +void WgRenderPipeline::allocate(WGPUDevice device, + WGPUVertexBufferLayout vertexBufferLayouts[], uint32_t attribsCount, + WGPUBindGroupLayout bindGroupLayouts[], uint32_t bindGroupsCount, + WGPUCompareFunction stencilCompareFunction, WGPUStencilOperation stencilOperation, + const char* shaderSource, const char* shaderLabel, const char* pipelineLabel) { mShaderModule = createShaderModule(device, shaderSource, shaderLabel); assert(mShaderModule); @@ -372,7 +440,8 @@ WGPUColorTargetState WgRenderPipeline::makeColorTargetState(const WGPUBlendState { WGPUColorTargetState colorTargetState{}; colorTargetState.nextInChain = nullptr; - colorTargetState.format = WGPUTextureFormat_BGRA8Unorm; // (WGPUTextureFormat_BGRA8UnormSrgb) + //colorTargetState.format = WGPUTextureFormat_BGRA8Unorm; // (WGPUTextureFormat_BGRA8UnormSrgb) + colorTargetState.format = WGPUTextureFormat_RGBA8Unorm; // (WGPUTextureFormat_BGRA8UnormSrgb) colorTargetState.blend = blendState; colorTargetState.writeMask = WGPUColorWriteMask_All; return colorTargetState; @@ -415,6 +484,7 @@ WGPUPrimitiveState WgRenderPipeline::makePrimitiveState() return primitiveState; } + WGPUDepthStencilState WgRenderPipeline::makeDepthStencilState(WGPUCompareFunction compare, WGPUStencilOperation operation) { WGPUDepthStencilState depthStencilState{}; @@ -463,11 +533,12 @@ WGPUFragmentState WgRenderPipeline::makeFragmentState(WGPUShaderModule shaderMod return fragmentState; } + WGPURenderPipeline WgRenderPipeline::createRenderPipeline(WGPUDevice device, - WGPUVertexBufferLayout vertexBufferLayouts[], uint32_t attribsCount, - WGPUCompareFunction stencilCompareFunction, WGPUStencilOperation stencilOperation, - WGPUPipelineLayout pipelineLayout, WGPUShaderModule shaderModule, - const char* pipelineName) + WGPUVertexBufferLayout vertexBufferLayouts[], uint32_t attribsCount, + WGPUCompareFunction stencilCompareFunction, WGPUStencilOperation stencilOperation, + WGPUPipelineLayout pipelineLayout, WGPUShaderModule shaderModule, + const char* pipelineName) { WGPUBlendState blendState = makeBlendState(); WGPUColorTargetState colorTargetStates[] = { @@ -492,8 +563,50 @@ WGPURenderPipeline WgRenderPipeline::createRenderPipeline(WGPUDevice device, return wgpuDeviceCreateRenderPipeline(device, &renderPipelineDesc); } + void WgRenderPipeline::destroyRenderPipeline(WGPURenderPipeline& renderPipeline) { if (renderPipeline) wgpuRenderPipelineRelease(renderPipeline); renderPipeline = nullptr; -} \ No newline at end of file +} + +//***************************************************************************** +// compute pipeline +//***************************************************************************** + +void WgComputePipeline::allocate(WGPUDevice device, + WGPUBindGroupLayout bindGroupLayouts[], uint32_t bindGroupsCount, + const char* shaderSource, const char* shaderLabel, const char* pipelineLabel) +{ + mShaderModule = createShaderModule(device, shaderSource, shaderLabel); + assert(mShaderModule); + + mPipelineLayout = createPipelineLayout(device, bindGroupLayouts, bindGroupsCount); + assert(mPipelineLayout); + + WGPUComputePipelineDescriptor computePipelineDesc{}; + computePipelineDesc.nextInChain = nullptr; + computePipelineDesc.label = pipelineLabel; + computePipelineDesc.layout = mPipelineLayout; + computePipelineDesc.compute.nextInChain = nullptr; + computePipelineDesc.compute.module = mShaderModule; + computePipelineDesc.compute.entryPoint = "cs_main"; + computePipelineDesc.compute.constantCount = 0; + computePipelineDesc.compute.constants = nullptr; + + mComputePipeline = wgpuDeviceCreateComputePipeline(device, &computePipelineDesc); + assert(mComputePipeline); +} + +void WgComputePipeline::release() +{ + if (mComputePipeline) wgpuComputePipelineRelease(mComputePipeline); + mComputePipeline = nullptr; + WgPipeline::release(); +} + + +void WgComputePipeline::set(WGPUComputePassEncoder computePassEncoder) +{ + wgpuComputePassEncoderSetPipeline(computePassEncoder, mComputePipeline); +} diff --git a/src/renderer/wg_engine/tvgWgCommon.h b/src/renderer/wg_engine/tvgWgCommon.h index 45349dce..5528eaa1 100644 --- a/src/renderer/wg_engine/tvgWgCommon.h +++ b/src/renderer/wg_engine/tvgWgCommon.h @@ -46,6 +46,11 @@ struct WgContext { void release(); void executeCommandEncoder(WGPUCommandEncoder commandEncoder); + + WGPUTexture createTexture2d(WGPUTextureUsage usage, WGPUTextureFormat format, uint32_t width, uint32_t height, char const * label); + WGPUTextureView createTextureView2d(WGPUTexture texture, WGPU_NULLABLE char const * label); + void releaseTexture(WGPUTexture& texture); + void releaseTextureView(WGPUTextureView& textureView); }; struct WgBindGroup @@ -53,6 +58,7 @@ struct WgBindGroup WGPUBindGroup mBindGroup{}; void set(WGPURenderPassEncoder encoder, uint32_t groupIndex); + void set(WGPUComputePassEncoder encoder, uint32_t groupIndex); static WGPUBindGroupEntry makeBindGroupEntryBuffer(uint32_t binding, WGPUBuffer buffer); static WGPUBindGroupEntry makeBindGroupEntrySampler(uint32_t binding, WGPUSampler sampler); @@ -60,7 +66,8 @@ struct WgBindGroup static WGPUBindGroupLayoutEntry makeBindGroupLayoutEntryBuffer(uint32_t binding); static WGPUBindGroupLayoutEntry makeBindGroupLayoutEntrySampler(uint32_t binding); - static WGPUBindGroupLayoutEntry makeBindGroupLayoutEntryTextureView(uint32_t binding); + static WGPUBindGroupLayoutEntry makeBindGroupLayoutEntryTexture(uint32_t binding); + static WGPUBindGroupLayoutEntry makeBindGroupLayoutEntryStorageTexture(uint32_t binding); 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); @@ -116,4 +123,17 @@ public: static void destroyRenderPipeline(WGPURenderPipeline& renderPipeline); }; +struct WgComputePipeline: public WgPipeline +{ +protected: + WGPUComputePipeline mComputePipeline{}; + void allocate(WGPUDevice device, + WGPUBindGroupLayout bindGroupLayouts[], uint32_t bindGroupsCount, + const char* shaderSource, const char* shaderLabel, const char* pipelineLabel); + +public: + void release() override; + void set(WGPUComputePassEncoder computePassEncoder); +}; + #endif // _TVG_WG_COMMON_H_ diff --git a/src/renderer/wg_engine/tvgWgPipelines.cpp b/src/renderer/wg_engine/tvgWgPipelines.cpp index 74bd9250..099b7319 100644 --- a/src/renderer/wg_engine/tvgWgPipelines.cpp +++ b/src/renderer/wg_engine/tvgWgPipelines.cpp @@ -235,7 +235,7 @@ void WgPipelineBlit::initialize(WGPUDevice device) // bind groups and layouts WGPUBindGroupLayout bindGroupLayouts[] = { - WgBindGroupBlit::getLayout(device), + WgBindGroupTextureSampled::getLayout(device), WgBindGroupOpacity::getLayout(device) }; @@ -269,7 +269,7 @@ void WgPipelineBlitColor::initialize(WGPUDevice device) // bind groups and layouts WGPUBindGroupLayout bindGroupLayouts[] = { - WgBindGroupBlit::getLayout(device) + WgBindGroupTextureSampled::getLayout(device) }; // stencil function @@ -302,8 +302,8 @@ void WgPipelineComposition::initialize(WGPUDevice device, const char* shaderSrc) // bind groups and layouts WGPUBindGroupLayout bindGroupLayouts[] = { - WgBindGroupBlit::getLayout(device), - WgBindGroupBlit::getLayout(device) + WgBindGroupTextureSampled::getLayout(device), + WgBindGroupTextureSampled::getLayout(device) }; // stencil function @@ -323,18 +323,40 @@ void WgPipelineComposition::initialize(WGPUDevice device, const char* shaderSrc) 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"; + + // allocate all pipeline handles + allocate(device, + bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts), + shaderSource, shaderLabel, pipelineLabel); +} + //************************************************************************ // pipelines //************************************************************************ void WgPipelines::initialize(WgContext& context) { + // fill pipelines fillShape.initialize(context.device); fillStroke.initialize(context.device); solid.initialize(context.device); linear.initialize(context.device); radial.initialize(context.device); image.initialize(context.device); + // blit pipelines blit.initialize(context.device); blitColor.initialize(context.device); // composition pipelines @@ -346,6 +368,8 @@ void WgPipelines::initialize(WgContext& context) 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); // store pipelines to context context.pipelines = this; } @@ -353,7 +377,9 @@ void WgPipelines::initialize(WgContext& context) void WgPipelines::release() { - WgBindGroupBlit::releaseLayout(); + WgBindGroupTextureSampled::releaseLayout(); + WgBindGroupStorageTexture::releaseLayout(); + WgBindGroupTexture::releaseLayout(); WgBindGroupOpacity::releaseLayout(); WgBindGroupPicture::releaseLayout(); WgBindGroupRadialGradient::releaseLayout(); @@ -361,6 +387,9 @@ void WgPipelines::release() WgBindGroupSolidColor::releaseLayout(); WgBindGroupPaint::releaseLayout(); WgBindGroupCanvas::releaseLayout(); + // compute pipelines + computeBlend.release(); + // composition pipelines compDifferenceMask.release(); compIntersectMask.release(); compSubtractMask.release(); @@ -369,8 +398,10 @@ void WgPipelines::release() compLumaMask.release(); compInvAlphaMask.release(); compAlphaMask.release(); + // blit pipelines blitColor.release(); blit.release(); + // fill pipelines image.release(); radial.release(); linear.release(); diff --git a/src/renderer/wg_engine/tvgWgPipelines.h b/src/renderer/wg_engine/tvgWgPipelines.h index ec60f8d8..8c34b23c 100644 --- a/src/renderer/wg_engine/tvgWgPipelines.h +++ b/src/renderer/wg_engine/tvgWgPipelines.h @@ -25,6 +25,10 @@ #include "tvgWgBindGroups.h" +//***************************************************************************** +// render pipelines +//***************************************************************************** + struct WgPipelineFillShape: public WgRenderPipeline { void initialize(WGPUDevice device) override; @@ -99,11 +103,11 @@ struct WgPipelineBlit: public WgRenderPipeline { void initialize(WGPUDevice device) override; void use(WGPURenderPassEncoder encoder, - WgBindGroupBlit& groupBlit, + WgBindGroupTextureSampled& groupTexSampled, WgBindGroupOpacity& groupOpacity) { set(encoder); - groupBlit.set(encoder, 0); + groupTexSampled.set(encoder, 0); groupOpacity.set(encoder, 1); } }; @@ -111,10 +115,11 @@ struct WgPipelineBlit: public WgRenderPipeline struct WgPipelineBlitColor: public WgRenderPipeline { void initialize(WGPUDevice device) override; - void use(WGPURenderPassEncoder encoder, WgBindGroupBlit& groupBlit) + void use(WGPURenderPassEncoder encoder, + WgBindGroupTextureSampled& groupTexSampled) { set(encoder); - groupBlit.set(encoder, 0); + groupTexSampled.set(encoder, 0); } }; @@ -122,14 +127,38 @@ struct WgPipelineComposition: public WgRenderPipeline { void initialize(WGPUDevice device) override {}; void initialize(WGPUDevice device, const char* shaderSrc); - void use(WGPURenderPassEncoder encoder, WgBindGroupBlit& groupBlitSrc, WgBindGroupBlit& groupBlitMsk) + void use(WGPURenderPassEncoder encoder, + WgBindGroupTextureSampled& groupTexSampledSrc, + WgBindGroupTextureSampled& groupTexSampledMsk) { set(encoder); - groupBlitSrc.set(encoder, 0); - groupBlitMsk.set(encoder, 1); + groupTexSampledSrc.set(encoder, 0); + groupTexSampledMsk.set(encoder, 1); } }; +//***************************************************************************** +// compute pipelines +//***************************************************************************** + +struct WgPipelineBlend: public WgComputePipeline +{ + void initialize(WGPUDevice device) override {}; + void initialize(WGPUDevice device, const char* shaderSrc); + void use(WGPUComputePassEncoder encoder, + WgBindGroupStorageTexture& groupTexSrc, + WgBindGroupStorageTexture& groupTexDst) + { + set(encoder); + groupTexSrc.set(encoder, 0); + groupTexDst.set(encoder, 1); + } +}; + +//***************************************************************************** +// pipelines +//***************************************************************************** + struct WgPipelines { WgPipelineFillShape fillShape; @@ -149,6 +178,8 @@ struct WgPipelines WgPipelineComposition compSubtractMask; WgPipelineComposition compIntersectMask; WgPipelineComposition compDifferenceMask; + // compute pipelines + WgPipelineBlend computeBlend; void initialize(WgContext& context); void release(); diff --git a/src/renderer/wg_engine/tvgWgRenderTarget.cpp b/src/renderer/wg_engine/tvgWgRenderTarget.cpp index 3f664488..193438ea 100644 --- a/src/renderer/wg_engine/tvgWgRenderTarget.cpp +++ b/src/renderer/wg_engine/tvgWgRenderTarget.cpp @@ -48,7 +48,7 @@ void WgRenderTarget::initialize(WgContext& context, uint32_t w, uint32_t h) textureDescColor.usage = WGPUTextureUsage_RenderAttachment | WGPUTextureUsage_TextureBinding | WGPUTextureUsage_CopyDst | WGPUTextureUsage_StorageBinding; textureDescColor.dimension = WGPUTextureDimension_2D; textureDescColor.size = { w, h, 1 }; - textureDescColor.format = WGPUTextureFormat_BGRA8Unorm; + textureDescColor.format = WGPUTextureFormat_RGBA8Unorm; textureDescColor.mipLevelCount = 1; textureDescColor.sampleCount = 1; textureDescColor.viewFormatCount = 0; @@ -59,7 +59,7 @@ void WgRenderTarget::initialize(WgContext& context, uint32_t w, uint32_t h) WGPUTextureViewDescriptor textureViewDescColor{}; textureViewDescColor.nextInChain = nullptr; textureViewDescColor.label = "The target texture view color"; - textureViewDescColor.format = WGPUTextureFormat_BGRA8Unorm; + textureViewDescColor.format = WGPUTextureFormat_RGBA8Unorm; textureViewDescColor.dimension = WGPUTextureViewDimension_2D; textureViewDescColor.baseMipLevel = 0; textureViewDescColor.mipLevelCount = 1; @@ -96,7 +96,9 @@ void WgRenderTarget::initialize(WgContext& context, uint32_t w, uint32_t h) textureViewStencil = wgpuTextureCreateView(mTextureStencil, &textureViewDescStencil); assert(textureViewStencil); // initialize bind group for blitting - bindGroupBlit.initialize(context.device, context.queue, sampler, textureViewColor); + bindGroupTex.initialize(context.device, context.queue, textureViewColor); + bindGroupStorageTex.initialize(context.device, context.queue, textureViewColor); + bindGroupTexSampled.initialize(context.device, context.queue, sampler, textureViewColor); // initialize window binding groups WgShaderTypeMat4x4f viewMat(w, h); mBindGroupCanvasWnd.initialize(context.device, context.queue, viewMat); @@ -111,7 +113,9 @@ void WgRenderTarget::release(WgContext& context) { mMeshDataCanvasWnd.release(context); mBindGroupCanvasWnd.release(); - bindGroupBlit.release(); + bindGroupTexSampled.release(); + bindGroupStorageTex.release(); + bindGroupTex.release(); if (mTextureStencil) { wgpuTextureDestroy(mTextureStencil); wgpuTextureRelease(mTextureStencil); @@ -167,7 +171,6 @@ void WgRenderTarget::beginRenderPass(WGPUCommandEncoder commandEncoder, WGPUText renderPassDesc.depthStencilAttachment = &depthStencilAttachment; //renderPassDesc.depthStencilAttachment = nullptr; renderPassDesc.occlusionQuerySet = nullptr; - renderPassDesc.timestampWriteCount = 0; renderPassDesc.timestampWrites = nullptr; // begin render pass mRenderPassEncoder = wgpuCommandEncoderBeginRenderPass(commandEncoder, &renderPassDesc); @@ -251,7 +254,7 @@ void WgRenderTarget::renderPicture(WgRenderDataPicture* renderData) void WgRenderTarget::blit(WgContext& context, WgRenderTarget* renderTargetSrc, WgBindGroupOpacity* mBindGroupOpacity) { assert(mRenderPassEncoder); - mPipelines->blit.use(mRenderPassEncoder, renderTargetSrc->bindGroupBlit, *mBindGroupOpacity); + mPipelines->blit.use(mRenderPassEncoder, renderTargetSrc->bindGroupTexSampled, *mBindGroupOpacity); mMeshDataCanvasWnd.drawImage(mRenderPassEncoder); } @@ -259,7 +262,7 @@ void WgRenderTarget::blit(WgContext& context, WgRenderTarget* renderTargetSrc, W void WgRenderTarget::blitColor(WgContext& context, WgRenderTarget* renderTargetSrc) { assert(mRenderPassEncoder); - mPipelines->blitColor.use(mRenderPassEncoder, renderTargetSrc->bindGroupBlit); + mPipelines->blitColor.use(mRenderPassEncoder, renderTargetSrc->bindGroupTexSampled); mMeshDataCanvasWnd.drawImage(mRenderPassEncoder); } @@ -269,7 +272,7 @@ void WgRenderTarget::compose(WgContext& context, WgRenderTarget* renderTargetSrc assert(mRenderPassEncoder); WgPipelineComposition* pipeline = mPipelines->getCompositionPipeline(method); assert(pipeline); - pipeline->use(mRenderPassEncoder, renderTargetSrc->bindGroupBlit, renderTargetMsk->bindGroupBlit); + pipeline->use(mRenderPassEncoder, renderTargetSrc->bindGroupTexSampled, renderTargetMsk->bindGroupTexSampled); mMeshDataCanvasWnd.drawImage(mRenderPassEncoder); } diff --git a/src/renderer/wg_engine/tvgWgRenderTarget.h b/src/renderer/wg_engine/tvgWgRenderTarget.h index e06d9997..f81b87d8 100644 --- a/src/renderer/wg_engine/tvgWgRenderTarget.h +++ b/src/renderer/wg_engine/tvgWgRenderTarget.h @@ -41,7 +41,9 @@ public: WGPUSampler sampler{}; WGPUTextureView textureViewColor{}; WGPUTextureView textureViewStencil{}; - WgBindGroupBlit bindGroupBlit; + WgBindGroupTexture bindGroupTex; + WgBindGroupStorageTexture bindGroupStorageTex; + WgBindGroupTextureSampled bindGroupTexSampled; public: void initialize(WgContext& context, uint32_t w, uint32_t h); void release(WgContext& context); diff --git a/src/renderer/wg_engine/tvgWgRenderer.cpp b/src/renderer/wg_engine/tvgWgRenderer.cpp index 21c14fd8..e6672e3d 100644 --- a/src/renderer/wg_engine/tvgWgRenderer.cpp +++ b/src/renderer/wg_engine/tvgWgRenderer.cpp @@ -43,6 +43,7 @@ void WgRenderer::initialize() { mContext.initialize(); mPipelines.initialize(mContext); + mBindGroupOpacityPool.initialize(mContext); } @@ -50,12 +51,12 @@ void WgRenderer::release() { mCompositorStack.clear(); mRenderTargetStack.clear(); - mBindGroupOpacityPool.release(); + mBindGroupOpacityPool.release(mContext); mRenderTargetPool.release(mContext); mRenderTargetRoot.release(mContext); mRenderTargetWnd.release(mContext); - if (mSwapChain) wgpuSwapChainRelease(mSwapChain); - if (mSurface) wgpuSurfaceRelease(mSurface); + wgpuSurfaceUnconfigure(mSurface); + wgpuSurfaceRelease(mSurface); mPipelines.release(); mContext.release(); } @@ -159,7 +160,6 @@ bool WgRenderer::postRender() { mRenderTargetRoot.endRenderPass(); mRenderTargetStack.pop(); - mBindGroupOpacityPool.reset(); mContext.executeCommandEncoder(mCommandEncoder); wgpuCommandEncoderRelease(mCommandEncoder); return true; @@ -211,7 +211,9 @@ bool WgRenderer::clear() bool WgRenderer::sync() { - WGPUTextureView backBufferView = wgpuSwapChainGetCurrentTextureView(mSwapChain); + WGPUSurfaceTexture backBuffer{}; + wgpuSurfaceGetCurrentTexture(mSurface, &backBuffer); + WGPUTextureView backBufferView = mContext.createTextureView2d(backBuffer.texture, "Surface texture view"); WGPUCommandEncoderDescriptor commandEncoderDesc{}; commandEncoderDesc.nextInChain = nullptr; commandEncoderDesc.label = "The command encoder"; @@ -222,7 +224,7 @@ bool WgRenderer::sync() mContext.executeCommandEncoder(commandEncoder); wgpuCommandEncoderRelease(commandEncoder); wgpuTextureViewRelease(backBufferView); - wgpuSwapChainPresent(mSwapChain); + wgpuSurfacePresent(mSurface); return true; } @@ -260,22 +262,22 @@ bool WgRenderer::target(void* window, uint32_t w, uint32_t h) mSurface = wgpuInstanceCreateSurface(mContext.instance, &surfaceDesc); assert(mSurface); - // get preferred format - WGPUTextureFormat swapChainFormat = WGPUTextureFormat_BGRA8Unorm; - // swapchain descriptor - WGPUSwapChainDescriptor swapChainDesc{}; - swapChainDesc.nextInChain = nullptr; - swapChainDesc.label = "The swapchain"; - swapChainDesc.usage = WGPUTextureUsage_RenderAttachment; - swapChainDesc.format = swapChainFormat; - swapChainDesc.width = mTargetSurface.w; - swapChainDesc.height = mTargetSurface.h; - swapChainDesc.presentMode = WGPUPresentMode_Mailbox; - mSwapChain = wgpuDeviceCreateSwapChain(mContext.device, mSurface, &swapChainDesc); - assert(mSwapChain); + WGPUSurfaceConfiguration surfaceConfiguration{}; + surfaceConfiguration.nextInChain = nullptr; + surfaceConfiguration.device = mContext.device; + surfaceConfiguration.format = WGPUTextureFormat_RGBA8Unorm; + surfaceConfiguration.usage = WGPUTextureUsage_RenderAttachment; + surfaceConfiguration.viewFormatCount = 0; + surfaceConfiguration.viewFormats = nullptr; + surfaceConfiguration.alphaMode = WGPUCompositeAlphaMode_Auto; + surfaceConfiguration.width = mTargetSurface.w; + surfaceConfiguration.height = mTargetSurface.h; + surfaceConfiguration.presentMode = WGPUPresentMode_Mailbox; + wgpuSurfaceConfigure(mSurface, &surfaceConfiguration); mRenderTargetWnd.initialize(mContext, w, h); mRenderTargetRoot.initialize(mContext, w, h); + return true; } diff --git a/src/renderer/wg_engine/tvgWgRenderer.h b/src/renderer/wg_engine/tvgWgRenderer.h index 004390b1..3def37f5 100644 --- a/src/renderer/wg_engine/tvgWgRenderer.h +++ b/src/renderer/wg_engine/tvgWgRenderer.h @@ -59,7 +59,7 @@ public: static WgRenderer* gen(); static bool init(uint32_t threads); static bool term(); - +private: // render handles WGPUCommandEncoder mCommandEncoder{}; Array mCompositorStack; @@ -74,7 +74,6 @@ private: WgRenderTarget mRenderTargetRoot; WgRenderTarget mRenderTargetWnd; WGPUSurface mSurface{}; - WGPUSwapChain mSwapChain{}; Surface mTargetSurface; }; diff --git a/src/renderer/wg_engine/tvgWgShaderSrc.cpp b/src/renderer/wg_engine/tvgWgShaderSrc.cpp index 74201fc7..49efb59c 100644 --- a/src/renderer/wg_engine/tvgWgShaderSrc.cpp +++ b/src/renderer/wg_engine/tvgWgShaderSrc.cpp @@ -707,3 +707,23 @@ 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; + +@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)); +} + +)"; diff --git a/src/renderer/wg_engine/tvgWgShaderSrc.h b/src/renderer/wg_engine/tvgWgShaderSrc.h index dfdebd26..448b85d2 100644 --- a/src/renderer/wg_engine/tvgWgShaderSrc.h +++ b/src/renderer/wg_engine/tvgWgShaderSrc.h @@ -25,26 +25,22 @@ #ifndef _TVG_WG_SHADER_SRC_H_ #define _TVG_WG_SHADER_SRC_H_ -// pipeline shader module fill +//***************************************************************************** +// render shader modules +//***************************************************************************** + +// pipeline shader modules fill extern const char* cShaderSource_PipelineFill; - -// pipeline shader module solid extern const char* cShaderSource_PipelineSolid; - -// pipeline shader module linear extern const char* cShaderSource_PipelineLinear; - -// pipeline shader module radial extern const char* cShaderSource_PipelineRadial; - -// pipeline shader module image extern const char* cShaderSource_PipelineImage; -// pipeline shader module blit +// pipeline shader modules blit extern const char* cShaderSource_PipelineBlit; extern const char* cShaderSource_PipelineBlitColor; -// pipeline shader module composes +// pipeline shader modules composes extern const char* cShaderSource_PipelineCompAlphaMask; extern const char* cShaderSource_PipelineCompInvAlphaMask; extern const char* cShaderSource_PipelineCompLumaMask; @@ -54,4 +50,11 @@ extern const char* cShaderSource_PipelineCompSubtractMask; extern const char* cShaderSource_PipelineCompIntersectMask; extern const char* cShaderSource_PipelineCompDifferenceMask; +//***************************************************************************** +// compute shader modules +//***************************************************************************** + +// pipeline shader modules blend +extern const char* cShaderSource_PipelineComputeBlend; + #endif // _TVG_WG_SHADER_SRC_H_