From f281cc9e92666a6ea2aed227695139a4a22fa4fa Mon Sep 17 00:00:00 2001 From: Sergii Liebodkin Date: Tue, 6 Aug 2024 19:47:49 +0000 Subject: [PATCH] wg_engine: composition and blend optimization * bind groups creation in real time removed - performance boost * blend and composition shaders decomposed - performance boost * shader modules and pipeline layouts generalized - less memory usage * shared single stencil buffer used - less memory usage * bind groups usage simplified * general context API simplified and generalized * all rendering logic moved into new composition class * ready for hardware MSAA (in next steps) * ready for direct mask applience (in next steps) --- src/renderer/wg_engine/meson.build | 2 + src/renderer/wg_engine/tvgWgBindGroups.cpp | 1047 ++++-------------- src/renderer/wg_engine/tvgWgBindGroups.h | 312 +----- src/renderer/wg_engine/tvgWgCommon.cpp | 827 +++----------- src/renderer/wg_engine/tvgWgCommon.h | 152 +-- src/renderer/wg_engine/tvgWgCompositor.cpp | 509 +++++++++ src/renderer/wg_engine/tvgWgCompositor.h | 90 ++ src/renderer/wg_engine/tvgWgGeometry.cpp | 0 src/renderer/wg_engine/tvgWgGeometry.h | 0 src/renderer/wg_engine/tvgWgPipelines.cpp | 747 +++++-------- src/renderer/wg_engine/tvgWgPipelines.h | 278 ++--- src/renderer/wg_engine/tvgWgRenderData.cpp | 93 +- src/renderer/wg_engine/tvgWgRenderData.h | 24 +- src/renderer/wg_engine/tvgWgRenderTarget.cpp | 426 +------ src/renderer/wg_engine/tvgWgRenderTarget.h | 100 +- src/renderer/wg_engine/tvgWgRenderer.cpp | 390 +++---- src/renderer/wg_engine/tvgWgRenderer.h | 37 +- src/renderer/wg_engine/tvgWgShaderSrc.cpp | 585 +++++----- src/renderer/wg_engine/tvgWgShaderSrc.h | 44 +- src/renderer/wg_engine/tvgWgShaderTypes.cpp | 0 src/renderer/wg_engine/tvgWgShaderTypes.h | 2 +- 21 files changed, 1949 insertions(+), 3716 deletions(-) mode change 100644 => 100755 src/renderer/wg_engine/meson.build mode change 100644 => 100755 src/renderer/wg_engine/tvgWgBindGroups.cpp mode change 100644 => 100755 src/renderer/wg_engine/tvgWgBindGroups.h mode change 100644 => 100755 src/renderer/wg_engine/tvgWgCommon.cpp mode change 100644 => 100755 src/renderer/wg_engine/tvgWgCommon.h create mode 100755 src/renderer/wg_engine/tvgWgCompositor.cpp create mode 100755 src/renderer/wg_engine/tvgWgCompositor.h mode change 100644 => 100755 src/renderer/wg_engine/tvgWgGeometry.cpp mode change 100644 => 100755 src/renderer/wg_engine/tvgWgGeometry.h mode change 100644 => 100755 src/renderer/wg_engine/tvgWgPipelines.cpp mode change 100644 => 100755 src/renderer/wg_engine/tvgWgPipelines.h mode change 100644 => 100755 src/renderer/wg_engine/tvgWgRenderData.cpp mode change 100644 => 100755 src/renderer/wg_engine/tvgWgRenderData.h mode change 100644 => 100755 src/renderer/wg_engine/tvgWgRenderTarget.cpp mode change 100644 => 100755 src/renderer/wg_engine/tvgWgRenderTarget.h mode change 100644 => 100755 src/renderer/wg_engine/tvgWgRenderer.cpp mode change 100644 => 100755 src/renderer/wg_engine/tvgWgRenderer.h mode change 100644 => 100755 src/renderer/wg_engine/tvgWgShaderSrc.cpp mode change 100644 => 100755 src/renderer/wg_engine/tvgWgShaderSrc.h mode change 100644 => 100755 src/renderer/wg_engine/tvgWgShaderTypes.cpp mode change 100644 => 100755 src/renderer/wg_engine/tvgWgShaderTypes.h diff --git a/src/renderer/wg_engine/meson.build b/src/renderer/wg_engine/meson.build old mode 100644 new mode 100755 index 95c0fe44..51666993 --- a/src/renderer/wg_engine/meson.build +++ b/src/renderer/wg_engine/meson.build @@ -1,6 +1,7 @@ source_file = [ 'tvgWgBindGroups.h', 'tvgWgCommon.h', + 'tvgWgCompositor.h', 'tvgWgGeometry.h', 'tvgWgPipelines.h', 'tvgWgRenderData.h', @@ -10,6 +11,7 @@ source_file = [ 'tvgWgShaderTypes.h', 'tvgWgBindGroups.cpp', 'tvgWgCommon.cpp', + 'tvgWgCompositor.cpp', 'tvgWgGeometry.cpp', 'tvgWgPipelines.cpp', 'tvgWgRenderData.cpp', diff --git a/src/renderer/wg_engine/tvgWgBindGroups.cpp b/src/renderer/wg_engine/tvgWgBindGroups.cpp old mode 100644 new mode 100755 index b051a6ab..012fe4ee --- a/src/renderer/wg_engine/tvgWgBindGroups.cpp +++ b/src/renderer/wg_engine/tvgWgBindGroups.cpp @@ -22,902 +22,225 @@ #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 groups -WGPUBindGroupLayout WgBindGroupTexture::layout = nullptr; -WGPUBindGroupLayout WgBindGroupTextureStorageRgbaRO::layout = nullptr; -WGPUBindGroupLayout WgBindGroupTextureStorageRgbaWO::layout = nullptr; -WGPUBindGroupLayout WgBindGroupTextureStorageBgraRO::layout = nullptr; -WGPUBindGroupLayout WgBindGroupTextureStorageBgraWO::layout = nullptr; -WGPUBindGroupLayout WgBindGroupTextureSampled::layout = nullptr; -WGPUBindGroupLayout WgBindGroupTexBlend::layout = nullptr; -WGPUBindGroupLayout WgBindGroupTexBlendMask::layout = nullptr; -WGPUBindGroupLayout WgBindGroupTexMaskCompose::layout = nullptr; -WGPUBindGroupLayout WgBindGroupTexCompose::layout = nullptr; -WGPUBindGroupLayout WgBindGroupOpacity::layout = nullptr; -WGPUBindGroupLayout WgBindGroupBlendMethod::layout = nullptr; -WGPUBindGroupLayout WgBindGroupCompositeMethod::layout = nullptr; - -/////////////////////////////////////////////////////////////////////////////// -// WgBindGroupCanvas -/////////////////////////////////////////////////////////////////////////////// - -WGPUBindGroupLayout WgBindGroupCanvas::getLayout(WGPUDevice device) +WGPUBindGroup WgBindGroupLayouts::createBindGroupTexSampled(WGPUSampler sampler, WGPUTextureView texView) { - if (layout) return layout; - const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { - makeBindGroupLayoutEntryBuffer(0) + const WGPUBindGroupEntry bindGroupEntrys[] = { + { .binding = 0, .sampler = sampler }, + { .binding = 1, .textureView = texView } }; - layout = createBindGroupLayout(device, bindGroupLayoutEntries, 1); - assert(layout); - return layout; + const WGPUBindGroupDescriptor bindGroupDesc { .layout = layoutTexSampled, .entryCount = 2, .entries = bindGroupEntrys }; + return wgpuDeviceCreateBindGroup(device, &bindGroupDesc); } -void WgBindGroupCanvas::releaseLayout() -{ - releaseBindGroupLayout(layout); -} - - -void WgBindGroupCanvas::initialize(WGPUDevice device, WGPUQueue queue, WgShaderTypeMat4x4f& uViewMat) -{ - release(); - uBufferViewMat = createBuffer(device, queue, &uViewMat, sizeof(uViewMat)); - const WGPUBindGroupEntry bindGroupEntries[] { - makeBindGroupEntryBuffer(0, uBufferViewMat) +WGPUBindGroup WgBindGroupLayouts::createBindGroupScreen1WO(WGPUTextureView texView) { + const WGPUBindGroupEntry bindGroupEntrys[] = { + { .binding = 0, .textureView = texView } }; - mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 1); - assert(mBindGroup); + const WGPUBindGroupDescriptor bindGroupDesc { .layout = layoutTexScreen1WO, .entryCount = 1, .entries = bindGroupEntrys }; + return wgpuDeviceCreateBindGroup(device, &bindGroupDesc); } -void WgBindGroupCanvas::release() +WGPUBindGroup WgBindGroupLayouts::createBindGroupStrorage1WO(WGPUTextureView texView) { - releaseBindGroup(mBindGroup); - releaseBuffer(uBufferViewMat); -} - -/////////////////////////////////////////////////////////////////////////////// -// WgBindGroupPaint -/////////////////////////////////////////////////////////////////////////////// - -WGPUBindGroupLayout WgBindGroupPaint::getLayout(WGPUDevice device) -{ - if (layout) return layout; - const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { - makeBindGroupLayoutEntryBuffer(0), - makeBindGroupLayoutEntryBuffer(1) + const WGPUBindGroupEntry bindGroupEntrys[] = { + { .binding = 0, .textureView = texView } }; - layout = createBindGroupLayout(device, bindGroupLayoutEntries, 2); - assert(layout); - return layout; + const WGPUBindGroupDescriptor bindGroupDesc { .layout = layoutTexStrorage1WO, .entryCount = 1, .entries = bindGroupEntrys }; + return wgpuDeviceCreateBindGroup(device, &bindGroupDesc); } -void WgBindGroupPaint::releaseLayout() +WGPUBindGroup WgBindGroupLayouts::createBindGroupStrorage1RO(WGPUTextureView texView) { - releaseBindGroupLayout(layout); + const WGPUBindGroupEntry bindGroupEntrys[] = { + { .binding = 0, .textureView = texView } + }; + const WGPUBindGroupDescriptor bindGroupDesc { .layout = layoutTexStrorage1RO, .entryCount = 1, .entries = bindGroupEntrys }; + return wgpuDeviceCreateBindGroup(device, &bindGroupDesc); } -void WgBindGroupPaint::initialize(WGPUDevice device, WGPUQueue queue, WgShaderTypeMat4x4f& uModelMat, WgShaderTypeBlendSettings& uBlendSettings) +WGPUBindGroup WgBindGroupLayouts::createBindGroupStrorage2RO(WGPUTextureView texView0, WGPUTextureView texView1) { - if (!uBufferModelMat && !uBufferBlendSettings && !mBindGroup) { - uBufferModelMat = createBuffer(device, queue, &uModelMat, sizeof(uModelMat)); - uBufferBlendSettings = createBuffer(device, queue, &uBlendSettings, sizeof(uBlendSettings)); - const WGPUBindGroupEntry bindGroupEntries[] { - makeBindGroupEntryBuffer(0, uBufferModelMat), - makeBindGroupEntryBuffer(1, uBufferBlendSettings) + const WGPUBindGroupEntry bindGroupEntrys[] = { + { .binding = 0, .textureView = texView0 }, + { .binding = 1, .textureView = texView1 } + }; + const WGPUBindGroupDescriptor bindGroupDesc { .layout = layoutTexStrorage2RO, .entryCount = 2, .entries = bindGroupEntrys }; + return wgpuDeviceCreateBindGroup(device, &bindGroupDesc); +} + + +WGPUBindGroup WgBindGroupLayouts::createBindGroupStrorage3RO(WGPUTextureView texView0, WGPUTextureView texView1, WGPUTextureView texView2) +{ + const WGPUBindGroupEntry bindGroupEntrys[] = { + { .binding = 0, .textureView = texView0 }, + { .binding = 1, .textureView = texView1 }, + { .binding = 2, .textureView = texView2 } + }; + const WGPUBindGroupDescriptor bindGroupDesc { .layout = layoutTexStrorage3RO, .entryCount = 3, .entries = bindGroupEntrys }; + return wgpuDeviceCreateBindGroup(device, &bindGroupDesc); +} + + +WGPUBindGroup WgBindGroupLayouts::createBindGroupBuffer1Un(WGPUBuffer buff) +{ + const WGPUBindGroupEntry bindGroupEntrys[] = { + { .binding = 0, .buffer = buff, .size = wgpuBufferGetSize(buff) } + }; + const WGPUBindGroupDescriptor bindGroupDesc { .layout = layoutBuffer1Un, .entryCount = 1, .entries = bindGroupEntrys }; + return wgpuDeviceCreateBindGroup(device, &bindGroupDesc); +} + + +WGPUBindGroup WgBindGroupLayouts::createBindGroupBuffer2Un(WGPUBuffer buff0, WGPUBuffer buff1) +{ + const WGPUBindGroupEntry bindGroupEntrys[] = { + { .binding = 0, .buffer = buff0, .size = wgpuBufferGetSize(buff0) }, + { .binding = 1, .buffer = buff1, .size = wgpuBufferGetSize(buff1) } + }; + const WGPUBindGroupDescriptor bindGroupDesc { .layout = layoutBuffer2Un, .entryCount = 2, .entries = bindGroupEntrys }; + return wgpuDeviceCreateBindGroup(device, &bindGroupDesc); +} + + +WGPUBindGroup WgBindGroupLayouts::createBindGroupBuffer3Un(WGPUBuffer buff0, WGPUBuffer buff1, WGPUBuffer buff2) +{ + const WGPUBindGroupEntry bindGroupEntrys[] = { + { .binding = 0, .buffer = buff0, .size = wgpuBufferGetSize(buff0) }, + { .binding = 1, .buffer = buff1, .size = wgpuBufferGetSize(buff1) }, + { .binding = 2, .buffer = buff2, .size = wgpuBufferGetSize(buff2) } + }; + const WGPUBindGroupDescriptor bindGroupDesc { .layout = layoutBuffer3Un, .entryCount = 3, .entries = bindGroupEntrys }; + return wgpuDeviceCreateBindGroup(device, &bindGroupDesc); +} + + +void WgBindGroupLayouts::releaseBindGroup(WGPUBindGroup& bindGroup) +{ + if (bindGroup) wgpuBindGroupRelease(bindGroup); + bindGroup = nullptr; +} + + +void WgBindGroupLayouts::initialize(WgContext& context) +{ + // store device handle + device = context.device; + assert(device); + + // common bind group settings + const WGPUShaderStageFlags visibility_vert = WGPUShaderStage_Vertex | WGPUShaderStage_Fragment | WGPUShaderStage_Compute; + const WGPUShaderStageFlags visibility_frag = WGPUShaderStage_Fragment | WGPUShaderStage_Compute; + const WGPUSamplerBindingLayout sampler = { .type = WGPUSamplerBindingType_Filtering }; + const WGPUTextureBindingLayout texture = { .sampleType = WGPUTextureSampleType_Float, .viewDimension = WGPUTextureViewDimension_2D }; + const WGPUStorageTextureBindingLayout storageTextureWO { .access = WGPUStorageTextureAccess_WriteOnly, .format = WGPUTextureFormat_RGBA8Unorm, .viewDimension = WGPUTextureViewDimension_2D }; + const WGPUStorageTextureBindingLayout storageTextureRO { .access = WGPUStorageTextureAccess_ReadOnly, .format = WGPUTextureFormat_RGBA8Unorm, .viewDimension = WGPUTextureViewDimension_2D }; + const WGPUStorageTextureBindingLayout storageScreenWO { .access = WGPUStorageTextureAccess_WriteOnly, .format = WGPUTextureFormat_BGRA8Unorm, .viewDimension = WGPUTextureViewDimension_2D }; + const WGPUBufferBindingLayout bufferUniform { .type = WGPUBufferBindingType_Uniform }; + + { // bind group layout tex sampled + const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { + { .binding = 0, .visibility = visibility_frag, .sampler = sampler }, + { .binding = 1, .visibility = visibility_frag, .texture = texture } }; - mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 2); - assert(mBindGroup); - return; + const WGPUBindGroupLayoutDescriptor bindGroupLayoutDesc { .entryCount = 2, .entries = bindGroupLayoutEntries }; + layoutTexSampled = wgpuDeviceCreateBindGroupLayout(device, &bindGroupLayoutDesc); + assert(layoutTexSampled); } - wgpuQueueWriteBuffer(queue, uBufferModelMat, 0, &uModelMat, sizeof(uModelMat)); - wgpuQueueWriteBuffer(queue, uBufferBlendSettings, 0, &uBlendSettings, sizeof(uBlendSettings)); -} - -void WgBindGroupPaint::release() -{ - releaseBindGroup(mBindGroup); - releaseBuffer(uBufferBlendSettings); - releaseBuffer(uBufferModelMat); -} - -/////////////////////////////////////////////////////////////////////////////// -// WgBindGroupSolidColor -/////////////////////////////////////////////////////////////////////////////// - -WGPUBindGroupLayout WgBindGroupSolidColor::getLayout(WGPUDevice device) -{ - if (layout) return layout; - const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { - makeBindGroupLayoutEntryBuffer(0) - }; - layout = createBindGroupLayout(device, bindGroupLayoutEntries, 1); - assert(layout); - return layout; -} - - -void WgBindGroupSolidColor::releaseLayout() -{ - releaseBindGroupLayout(layout); -} - - -void WgBindGroupSolidColor::initialize(WGPUDevice device, WGPUQueue queue, WgShaderTypeSolidColor &uSolidColor) -{ - if (!uBufferSolidColor && !mBindGroup) { - uBufferSolidColor = createBuffer(device, queue, &uSolidColor, sizeof(uSolidColor)); - const WGPUBindGroupEntry bindGroupEntries[] { - makeBindGroupEntryBuffer(0, uBufferSolidColor) + { // bind group layout tex screen 1 RO + const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { + { .binding = 0, .visibility = visibility_frag, .storageTexture = storageScreenWO } }; - mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 1); - assert(mBindGroup); - return; + const WGPUBindGroupLayoutDescriptor bindGroupLayoutDesc { .entryCount = 1, .entries = bindGroupLayoutEntries }; + layoutTexScreen1WO = wgpuDeviceCreateBindGroupLayout(device, &bindGroupLayoutDesc); + assert(layoutTexScreen1WO); } - wgpuQueueWriteBuffer(queue, uBufferSolidColor, 0, &uSolidColor, sizeof(uSolidColor)); -} - -void WgBindGroupSolidColor::release() -{ - releaseBindGroup(mBindGroup); - releaseBuffer(uBufferSolidColor); -} - -/////////////////////////////////////////////////////////////////////////////// -// WgBindGroupLinearGradient -/////////////////////////////////////////////////////////////////////////////// - -WGPUBindGroupLayout WgBindGroupLinearGradient::getLayout(WGPUDevice device) -{ - if (layout) return layout; - const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { - makeBindGroupLayoutEntryBuffer(0) - }; - layout = createBindGroupLayout(device, bindGroupLayoutEntries, 1); - assert(layout); - return layout; -} - - -void WgBindGroupLinearGradient::releaseLayout() -{ - releaseBindGroupLayout(layout); -} - - -void WgBindGroupLinearGradient::initialize(WGPUDevice device, WGPUQueue queue, WgShaderTypeLinearGradient &uLinearGradient) -{ - if (!uBufferLinearGradient && !mBindGroup) { - uBufferLinearGradient = createBuffer(device, queue, &uLinearGradient, sizeof(uLinearGradient)); - const WGPUBindGroupEntry bindGroupEntries[] { - makeBindGroupEntryBuffer(0, uBufferLinearGradient) + { // bind group layout tex storage 1 WO + const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { + { .binding = 0, .visibility = visibility_frag, .storageTexture = storageTextureWO } }; - mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 1); - assert(mBindGroup); - return; + const WGPUBindGroupLayoutDescriptor bindGroupLayoutDesc { .entryCount = 1, .entries = bindGroupLayoutEntries }; + layoutTexStrorage1WO = wgpuDeviceCreateBindGroupLayout(device, &bindGroupLayoutDesc); + assert(layoutTexStrorage1WO); } - wgpuQueueWriteBuffer(queue, uBufferLinearGradient, 0, &uLinearGradient, sizeof(uLinearGradient)); -} - -void WgBindGroupLinearGradient::release() -{ - releaseBindGroup(mBindGroup); - releaseBuffer(uBufferLinearGradient); -} - -/////////////////////////////////////////////////////////////////////////////// -// WgBindGroupRadialGradient -/////////////////////////////////////////////////////////////////////////////// - -WGPUBindGroupLayout WgBindGroupRadialGradient::getLayout(WGPUDevice device) -{ - if (layout) return layout; - const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { - makeBindGroupLayoutEntryBuffer(0) - }; - layout = createBindGroupLayout(device, bindGroupLayoutEntries, 1); - assert(layout); - return layout; -} - - -void WgBindGroupRadialGradient::releaseLayout() -{ - releaseBindGroupLayout(layout); -} - - -void WgBindGroupRadialGradient::initialize(WGPUDevice device, WGPUQueue queue, WgShaderTypeRadialGradient &uRadialGradient) -{ - if (!uBufferRadialGradient && !mBindGroup) { - uBufferRadialGradient = createBuffer(device, queue, &uRadialGradient, sizeof(uRadialGradient)); - const WGPUBindGroupEntry bindGroupEntries[] { - makeBindGroupEntryBuffer(0, uBufferRadialGradient) + { // bind group layout tex storage 1 RO + const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { + { .binding = 0, .visibility = visibility_frag, .storageTexture = storageTextureRO } }; - mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 1); - assert(mBindGroup); - return; + const WGPUBindGroupLayoutDescriptor bindGroupLayoutDesc { .entryCount = 1, .entries = bindGroupLayoutEntries }; + layoutTexStrorage1RO = wgpuDeviceCreateBindGroupLayout(device, &bindGroupLayoutDesc); + assert(layoutTexStrorage1RO); } - wgpuQueueWriteBuffer(queue, uBufferRadialGradient, 0, &uRadialGradient, sizeof(uRadialGradient)); -} - -void WgBindGroupRadialGradient::release() -{ - releaseBuffer(uBufferRadialGradient); - releaseBindGroup(mBindGroup); -} - -/////////////////////////////////////////////////////////////////////////////// -// WgBindGroupPicture -/////////////////////////////////////////////////////////////////////////////// - -WGPUBindGroupLayout WgBindGroupPicture::getLayout(WGPUDevice device) -{ - if (layout) return layout; - const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { - makeBindGroupLayoutEntrySampler(0), - makeBindGroupLayoutEntryTexture(1) - }; - layout = createBindGroupLayout(device, bindGroupLayoutEntries, 2); - assert(layout); - return layout; -} - - -void WgBindGroupPicture::releaseLayout() -{ - releaseBindGroupLayout(layout); -} - - -void WgBindGroupPicture::initialize(WGPUDevice device, WGPUQueue queue, WGPUSampler uSampler, WGPUTextureView uTextureView) -{ - release(); - const WGPUBindGroupEntry bindGroupEntries[] { - makeBindGroupEntrySampler(0, uSampler), - makeBindGroupEntryTextureView(1, uTextureView) - }; - mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 2); - assert(mBindGroup); -} - - -void WgBindGroupPicture::release() -{ - releaseBindGroup(mBindGroup); -} - -/////////////////////////////////////////////////////////////////////////////// -// WgBindGroupTexture -/////////////////////////////////////////////////////////////////////////////// - -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); -} - -/////////////////////////////////////////////////////////////////////////////// -// WgBindGroupTextureStorageRgbaRO -/////////////////////////////////////////////////////////////////////////////// - -WGPUBindGroupLayout WgBindGroupTextureStorageRgbaRO::getLayout(WGPUDevice device) -{ - if (layout) return layout; - const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { - makeBindGroupLayoutEntryStorage(0, WGPUStorageTextureAccess_ReadOnly, WGPUTextureFormat_RGBA8Unorm) - }; - layout = createBindGroupLayout(device, bindGroupLayoutEntries, 1); - assert(layout); - return layout; -} - - -void WgBindGroupTextureStorageRgbaRO::releaseLayout() -{ - releaseBindGroupLayout(layout); -} - - -void WgBindGroupTextureStorageRgbaRO::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexture) -{ - release(); - const WGPUBindGroupEntry bindGroupEntries[] { - makeBindGroupEntryTextureView(0, uTexture) - }; - mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 1); - assert(mBindGroup); -} - - -void WgBindGroupTextureStorageRgbaRO::release() -{ - releaseBindGroup(mBindGroup); -} - -/////////////////////////////////////////////////////////////////////////////// -// WgBindGroupTextureStorageRgbaRO -/////////////////////////////////////////////////////////////////////////////// - -WGPUBindGroupLayout WgBindGroupTextureStorageRgbaWO::getLayout(WGPUDevice device) -{ - if (layout) return layout; - const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { - makeBindGroupLayoutEntryStorage(0, WGPUStorageTextureAccess_WriteOnly, WGPUTextureFormat_RGBA8Unorm) - }; - layout = createBindGroupLayout(device, bindGroupLayoutEntries, 1); - assert(layout); - return layout; -} - - -void WgBindGroupTextureStorageRgbaWO::releaseLayout() -{ - releaseBindGroupLayout(layout); -} - - -void WgBindGroupTextureStorageRgbaWO::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexture) -{ - release(); - const WGPUBindGroupEntry bindGroupEntries[] { - makeBindGroupEntryTextureView(0, uTexture) - }; - mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 1); - assert(mBindGroup); -} - -void WgBindGroupTextureStorageRgbaWO::release() -{ - releaseBindGroup(mBindGroup); -} - -/////////////////////////////////////////////////////////////////////////////// -// WgBindGroupTextureStorageBgraRO -/////////////////////////////////////////////////////////////////////////////// - -WGPUBindGroupLayout WgBindGroupTextureStorageBgraRO::getLayout(WGPUDevice device) -{ - if (layout) return layout; - const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { - makeBindGroupLayoutEntryStorage(0, WGPUStorageTextureAccess_ReadOnly, WGPUTextureFormat_BGRA8Unorm) - }; - layout = createBindGroupLayout(device, bindGroupLayoutEntries, 1); - assert(layout); - return layout; -} - - -void WgBindGroupTextureStorageBgraRO::releaseLayout() -{ - releaseBindGroupLayout(layout); -} - - -void WgBindGroupTextureStorageBgraRO::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexture) -{ - release(); - const WGPUBindGroupEntry bindGroupEntries[] { - makeBindGroupEntryTextureView(0, uTexture) - }; - mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 1); - assert(mBindGroup); -} - - -void WgBindGroupTextureStorageBgraRO::release() -{ - releaseBindGroup(mBindGroup); -} - -/////////////////////////////////////////////////////////////////////////////// -// WgBindGroupTextureStorageBgraWO -/////////////////////////////////////////////////////////////////////////////// - -WGPUBindGroupLayout WgBindGroupTextureStorageBgraWO::getLayout(WGPUDevice device) -{ - if (layout) return layout; - const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { - makeBindGroupLayoutEntryStorage(0, WGPUStorageTextureAccess_WriteOnly, WGPUTextureFormat_BGRA8Unorm) - }; - layout = createBindGroupLayout(device, bindGroupLayoutEntries, 1); - assert(layout); - return layout; -} - - -void WgBindGroupTextureStorageBgraWO::releaseLayout() -{ - releaseBindGroupLayout(layout); -} - - -void WgBindGroupTextureStorageBgraWO::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexture) -{ - release(); - const WGPUBindGroupEntry bindGroupEntries[] { - makeBindGroupEntryTextureView(0, uTexture) - }; - mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 1); - assert(mBindGroup); -} - - -void WgBindGroupTextureStorageBgraWO::release() -{ - releaseBindGroup(mBindGroup); -} - -/////////////////////////////////////////////////////////////////////////////// -// WgBindGroupTextureSampled -/////////////////////////////////////////////////////////////////////////////// - -WGPUBindGroupLayout WgBindGroupTextureSampled::getLayout(WGPUDevice device) -{ - if (layout) return layout; - const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { - makeBindGroupLayoutEntrySampler(0), - makeBindGroupLayoutEntryTexture(1) - }; - layout = createBindGroupLayout(device, bindGroupLayoutEntries, 2); - assert(layout); - return layout; -} - - -void WgBindGroupTextureSampled::releaseLayout() -{ - releaseBindGroupLayout(layout); -} - - -void WgBindGroupTextureSampled::initialize(WGPUDevice device, WGPUQueue queue, WGPUSampler uSampler, WGPUTextureView uTexture) -{ - release(); - const WGPUBindGroupEntry bindGroupEntries[] { - makeBindGroupEntrySampler(0, uSampler), - makeBindGroupEntryTextureView(1, uTexture) - }; - mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 2); - assert(mBindGroup); -} - - -void WgBindGroupTextureSampled::release() -{ - releaseBindGroup(mBindGroup); -} - -/////////////////////////////////////////////////////////////////////////////// -// WgBindGroupTexBlend -/////////////////////////////////////////////////////////////////////////////// - -WGPUBindGroupLayout WgBindGroupTexBlend::getLayout(WGPUDevice device) -{ - if (layout) return layout; - const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { - makeBindGroupLayoutEntryStorage(0, WGPUStorageTextureAccess_ReadOnly, WGPUTextureFormat_RGBA8Unorm), - makeBindGroupLayoutEntryStorage(1, WGPUStorageTextureAccess_ReadOnly, WGPUTextureFormat_RGBA8Unorm), - makeBindGroupLayoutEntryStorage(2, WGPUStorageTextureAccess_WriteOnly, WGPUTextureFormat_RGBA8Unorm) - }; - layout = createBindGroupLayout(device, bindGroupLayoutEntries, 3); - assert(layout); - return layout; -} - - -void WgBindGroupTexBlend::releaseLayout() -{ - releaseBindGroupLayout(layout); -} - - -void WgBindGroupTexBlend::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexSrc, WGPUTextureView uTexMsk, WGPUTextureView uTexDst) -{ - release(); - const WGPUBindGroupEntry bindGroupEntries[] { - makeBindGroupEntryTextureView(0, uTexSrc), - makeBindGroupEntryTextureView(1, uTexMsk), - makeBindGroupEntryTextureView(2, uTexDst) - }; - mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 3); - assert(mBindGroup); -} - - -void WgBindGroupTexBlend::release() -{ - releaseBindGroup(mBindGroup); -} - -/////////////////////////////////////////////////////////////////////////////// -// WgBindGroupTexBlendMask -/////////////////////////////////////////////////////////////////////////////// - -WGPUBindGroupLayout WgBindGroupTexBlendMask::getLayout(WGPUDevice device) -{ - if (layout) return layout; - const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { - makeBindGroupLayoutEntryStorage(0, WGPUStorageTextureAccess_ReadOnly, WGPUTextureFormat_RGBA8Unorm), - makeBindGroupLayoutEntryStorage(1, WGPUStorageTextureAccess_ReadOnly, WGPUTextureFormat_RGBA8Unorm), - makeBindGroupLayoutEntryStorage(2, WGPUStorageTextureAccess_ReadOnly, WGPUTextureFormat_RGBA8Unorm), - makeBindGroupLayoutEntryStorage(3, WGPUStorageTextureAccess_WriteOnly, WGPUTextureFormat_RGBA8Unorm) - }; - layout = createBindGroupLayout(device, bindGroupLayoutEntries, 4); - assert(layout); - return layout; -} - - -void WgBindGroupTexBlendMask::releaseLayout() -{ - releaseBindGroupLayout(layout); -} - - -void WgBindGroupTexBlendMask::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexSrc, WGPUTextureView uTexMsk, WGPUTextureView uTexDst, WGPUTextureView uTexTrg) -{ - release(); - const WGPUBindGroupEntry bindGroupEntries[] { - makeBindGroupEntryTextureView(0, uTexSrc), - makeBindGroupEntryTextureView(1, uTexMsk), - makeBindGroupEntryTextureView(2, uTexDst), - makeBindGroupEntryTextureView(3, uTexTrg) - }; - mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 4); - assert(mBindGroup); -} - - -void WgBindGroupTexBlendMask::release() -{ - releaseBindGroup(mBindGroup); -} - -/////////////////////////////////////////////////////////////////////////////// -// WgBindGroupTexMaskCompose -/////////////////////////////////////////////////////////////////////////////// - -WGPUBindGroupLayout WgBindGroupTexMaskCompose::getLayout(WGPUDevice device) -{ - if (layout) return layout; - const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { - makeBindGroupLayoutEntryStorage(0, WGPUStorageTextureAccess_ReadOnly, WGPUTextureFormat_RGBA8Unorm), - makeBindGroupLayoutEntryStorage(1, WGPUStorageTextureAccess_ReadOnly, WGPUTextureFormat_RGBA8Unorm), - makeBindGroupLayoutEntryStorage(2, WGPUStorageTextureAccess_WriteOnly, WGPUTextureFormat_RGBA8Unorm) - }; - layout = createBindGroupLayout(device, bindGroupLayoutEntries, 3); - assert(layout); - return layout; -} - - -void WgBindGroupTexMaskCompose::releaseLayout() -{ - releaseBindGroupLayout(layout); -} - - -void WgBindGroupTexMaskCompose::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexMsk0, WGPUTextureView uTexMsk1, WGPUTextureView uTexTrg) -{ - release(); - const WGPUBindGroupEntry bindGroupEntries[] { - makeBindGroupEntryTextureView(0, uTexMsk0), - makeBindGroupEntryTextureView(1, uTexMsk1), - makeBindGroupEntryTextureView(2, uTexTrg) - }; - mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 3); - assert(mBindGroup); -} - - -void WgBindGroupTexMaskCompose::release() -{ - releaseBindGroup(mBindGroup); -} - -/////////////////////////////////////////////////////////////////////////////// -// WgBindGroupTexCompose -/////////////////////////////////////////////////////////////////////////////// - -WGPUBindGroupLayout WgBindGroupTexCompose::getLayout(WGPUDevice device) -{ - if (layout) return layout; - const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { - makeBindGroupLayoutEntryStorage(0, WGPUStorageTextureAccess_ReadOnly, WGPUTextureFormat_RGBA8Unorm), - makeBindGroupLayoutEntryStorage(1, WGPUStorageTextureAccess_ReadOnly, WGPUTextureFormat_RGBA8Unorm), - makeBindGroupLayoutEntryStorage(2, WGPUStorageTextureAccess_ReadOnly, WGPUTextureFormat_RGBA8Unorm), - makeBindGroupLayoutEntryStorage(3, WGPUStorageTextureAccess_WriteOnly, WGPUTextureFormat_RGBA8Unorm) - }; - layout = createBindGroupLayout(device, bindGroupLayoutEntries, 4); - assert(layout); - return layout; -} - - -void WgBindGroupTexCompose::releaseLayout() -{ - releaseBindGroupLayout(layout); -} - - -void WgBindGroupTexCompose::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexSrc, WGPUTextureView uTexMsk, WGPUTextureView uTexDst, WGPUTextureView uTexTrg) -{ - release(); - const WGPUBindGroupEntry bindGroupEntries[] { - makeBindGroupEntryTextureView(0, uTexSrc), - makeBindGroupEntryTextureView(1, uTexMsk), - makeBindGroupEntryTextureView(2, uTexDst), - makeBindGroupEntryTextureView(3, uTexTrg) - }; - mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 4); - assert(mBindGroup); -} - - -void WgBindGroupTexCompose::release() -{ - releaseBindGroup(mBindGroup); -} - -/////////////////////////////////////////////////////////////////////////////// -// WgBindGroupOpacity -/////////////////////////////////////////////////////////////////////////////// - -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); -} - -/////////////////////////////////////////////////////////////////////////////// -// WgBindGroupBlendMethod -/////////////////////////////////////////////////////////////////////////////// - -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); -} - -/////////////////////////////////////////////////////////////////////////////// -// WgBindGroupCompositeMethod -/////////////////////////////////////////////////////////////////////////////// - -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) -{ - release(context); - for (uint32_t opacity = 0; opacity < 256; opacity++) { - mPool[opacity] = new WgBindGroupOpacity; - mPool[opacity]->initialize(context.device, context.queue, opacity); + { // bind group layout tex storage 2 RO + const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { + { .binding = 0, .visibility = visibility_frag, .storageTexture = storageTextureRO }, + { .binding = 1, .visibility = visibility_frag, .storageTexture = storageTextureRO } + }; + const WGPUBindGroupLayoutDescriptor bindGroupLayoutDesc { .entryCount = 2, .entries = bindGroupLayoutEntries }; + layoutTexStrorage2RO = wgpuDeviceCreateBindGroupLayout(device, &bindGroupLayoutDesc); + assert(layoutTexStrorage2RO); + } + + { // bind group layout tex storage 3 RO + const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { + { .binding = 0, .visibility = visibility_frag, .storageTexture = storageTextureRO }, + { .binding = 1, .visibility = visibility_frag, .storageTexture = storageTextureRO }, + { .binding = 2, .visibility = visibility_frag, .storageTexture = storageTextureRO } + }; + const WGPUBindGroupLayoutDescriptor bindGroupLayoutDesc { .entryCount = 3, .entries = bindGroupLayoutEntries }; + layoutTexStrorage3RO = wgpuDeviceCreateBindGroupLayout(device, &bindGroupLayoutDesc); + assert(layoutTexStrorage3RO); + } + + { // bind group layout buffer 1 uniform + const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { + { .binding = 0, .visibility = visibility_vert, .buffer = bufferUniform } + }; + const WGPUBindGroupLayoutDescriptor bindGroupLayoutDesc { .entryCount = 1, .entries = bindGroupLayoutEntries }; + layoutBuffer1Un = wgpuDeviceCreateBindGroupLayout(device, &bindGroupLayoutDesc); + assert(layoutBuffer1Un); + } + + { // bind group layout buffer 2 uniform + const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { + { .binding = 0, .visibility = visibility_vert, .buffer = bufferUniform }, + { .binding = 1, .visibility = visibility_vert, .buffer = bufferUniform } + }; + const WGPUBindGroupLayoutDescriptor bindGroupLayoutDesc { .entryCount = 2, .entries = bindGroupLayoutEntries }; + layoutBuffer2Un = wgpuDeviceCreateBindGroupLayout(device, &bindGroupLayoutDesc); + assert(layoutBuffer2Un); + } + + { // bind group layout buffer 3 uniform + const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { + { .binding = 0, .visibility = visibility_vert, .buffer = bufferUniform }, + { .binding = 1, .visibility = visibility_vert, .buffer = bufferUniform }, + { .binding = 2, .visibility = visibility_vert, .buffer = bufferUniform } + }; + const WGPUBindGroupLayoutDescriptor bindGroupLayoutDesc { .entryCount = 3, .entries = bindGroupLayoutEntries }; + layoutBuffer3Un = wgpuDeviceCreateBindGroupLayout(device, &bindGroupLayoutDesc); + assert(layoutBuffer3Un); } } -void WgBindGroupOpacityPool::release(WgContext& context) +void WgBindGroupLayouts::release(WgContext& context) { - for (uint32_t i = 0; i < 256; i++) { - if (mPool[i]) { - mPool[i]->release(); - delete mPool[i]; - mPool[i] = nullptr; - } - } -} - - -WgBindGroupOpacity* WgBindGroupOpacityPool::allocate(WgContext& context, uint8_t opacity) -{ - 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 composeMethods = (uint8_t)CompositeMethod::None; - composeMethods <= (uint8_t)CompositeMethod::DifferenceMask; - composeMethods++) { - mPool[composeMethods] = new WgBindGroupCompositeMethod; - mPool[composeMethods]->initialize(context.device, context.queue, (CompositeMethod)composeMethods); - } -} - - -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]; + wgpuBindGroupLayoutRelease(layoutBuffer3Un); + wgpuBindGroupLayoutRelease(layoutBuffer2Un); + wgpuBindGroupLayoutRelease(layoutBuffer1Un); + wgpuBindGroupLayoutRelease(layoutTexScreen1WO); + wgpuBindGroupLayoutRelease(layoutTexStrorage3RO); + wgpuBindGroupLayoutRelease(layoutTexStrorage2RO); + wgpuBindGroupLayoutRelease(layoutTexStrorage1RO); + wgpuBindGroupLayoutRelease(layoutTexStrorage1WO); + wgpuBindGroupLayoutRelease(layoutTexSampled); + device = nullptr; } diff --git a/src/renderer/wg_engine/tvgWgBindGroups.h b/src/renderer/wg_engine/tvgWgBindGroups.h old mode 100644 new mode 100755 index d1837b91..6780442e --- a/src/renderer/wg_engine/tvgWgBindGroups.h +++ b/src/renderer/wg_engine/tvgWgBindGroups.h @@ -24,300 +24,34 @@ #define _TVG_WG_BIND_GROUPS_H_ #include "tvgWgCommon.h" -#include "tvgWgShaderTypes.h" -// @group(0) -struct WgBindGroupCanvas : public WgBindGroup -{ - static WGPUBindGroupLayout layout; - static WGPUBindGroupLayout getLayout(WGPUDevice device); - static void releaseLayout(); - - WGPUBuffer uBufferViewMat{}; - void initialize(WGPUDevice device, WGPUQueue queue, - WgShaderTypeMat4x4f& uViewMat); - void release(); -}; - -// @group(1) -struct WgBindGroupPaint : public WgBindGroup -{ - static WGPUBindGroupLayout layout; - static WGPUBindGroupLayout getLayout(WGPUDevice device); - static void releaseLayout(); - - WGPUBuffer uBufferModelMat{}; - WGPUBuffer uBufferBlendSettings{}; - void initialize(WGPUDevice device, WGPUQueue queue, - WgShaderTypeMat4x4f& uModelMat, - WgShaderTypeBlendSettings& uBlendSettings); - void release(); -}; - -// @group(2) -struct WgBindGroupSolidColor : public WgBindGroup -{ - static WGPUBindGroupLayout layout; - static WGPUBindGroupLayout getLayout(WGPUDevice device); - static void releaseLayout(); - - WGPUBuffer uBufferSolidColor{}; - void initialize(WGPUDevice device, WGPUQueue queue, - WgShaderTypeSolidColor &uSolidColor); - void release(); -}; - -// @group(2) -struct WgBindGroupLinearGradient : public WgBindGroup -{ - static WGPUBindGroupLayout layout; - static WGPUBindGroupLayout getLayout(WGPUDevice device); - static void releaseLayout(); - - WGPUBuffer uBufferLinearGradient{}; - void initialize(WGPUDevice device, WGPUQueue queue, - WgShaderTypeLinearGradient &uLinearGradient); - void release(); -}; - -// @group(2) -struct WgBindGroupRadialGradient : public WgBindGroup -{ - static WGPUBindGroupLayout layout; - static WGPUBindGroupLayout getLayout(WGPUDevice device); - static void releaseLayout(); - - WGPUBuffer uBufferRadialGradient{}; - void initialize(WGPUDevice device, WGPUQueue queue, - WgShaderTypeRadialGradient &uRadialGradient); - void release(); -}; - -// @group(2) -struct WgBindGroupPicture : public WgBindGroup -{ - static WGPUBindGroupLayout layout; - static WGPUBindGroupLayout getLayout(WGPUDevice device); - static void releaseLayout(); - - void initialize(WGPUDevice device, WGPUQueue queue, - WGPUSampler uSampler, - WGPUTextureView uTextureView); - void release(); -}; - -// @group(0 or 1) -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 WgBindGroupTextureStorageRgbaWO : 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 WgBindGroupTextureStorageRgbaRO : 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 WgBindGroupTextureStorageBgraRO : public WgBindGroup -{ - static WGPUBindGroupLayout layout; - static WGPUBindGroupLayout getLayout(WGPUDevice device); - static WGPUBindGroupLayout getLayoutRO(WGPUDevice device); - static WGPUBindGroupLayout getLayoutWO(WGPUDevice device); - static void releaseLayout(); - - void initialize(WGPUDevice device, WGPUQueue queue, - WGPUTextureView uTexture); - void release(); -}; - - -// @group(0 or 1) -struct WgBindGroupTextureStorageBgraWO : public WgBindGroup -{ - static WGPUBindGroupLayout layout; - static WGPUBindGroupLayout getLayout(WGPUDevice device); - static WGPUBindGroupLayout getLayoutRO(WGPUDevice device); - static WGPUBindGroupLayout getLayoutWO(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); - static void releaseLayout(); - - void initialize(WGPUDevice device, WGPUQueue queue, - WGPUSampler uSampler, - WGPUTextureView uTexture); - void release(); -}; - -// @group(0) -struct WgBindGroupTexBlend : public WgBindGroup -{ - static WGPUBindGroupLayout layout; - static WGPUBindGroupLayout getLayout(WGPUDevice device); - static void releaseLayout(); - - void initialize(WGPUDevice device, WGPUQueue queue, - WGPUTextureView uTexSrc, - WGPUTextureView uTexDst, - WGPUTextureView uTexTrg); - void release(); -}; - - -// @group(0) -struct WgBindGroupTexBlendMask : public WgBindGroup -{ - static WGPUBindGroupLayout layout; - static WGPUBindGroupLayout getLayout(WGPUDevice device); - static void releaseLayout(); - - void initialize(WGPUDevice device, WGPUQueue queue, - WGPUTextureView uTexSrc, - WGPUTextureView uTexMsk, - WGPUTextureView uTexDst, - WGPUTextureView uTexTrg); - void release(); -}; - - -// @group(0) -struct WgBindGroupTexMaskCompose : public WgBindGroup -{ - static WGPUBindGroupLayout layout; - static WGPUBindGroupLayout getLayout(WGPUDevice device); - static void releaseLayout(); - - void initialize(WGPUDevice device, WGPUQueue queue, - WGPUTextureView uTexMsk0, - WGPUTextureView uTexMsk1, - WGPUTextureView uTexTrg); - void release(); -}; - - -// @group(0) -struct WgBindGroupTexCompose : public WgBindGroup -{ - static WGPUBindGroupLayout layout; - static WGPUBindGroupLayout getLayout(WGPUDevice device); - static void releaseLayout(); - - void initialize(WGPUDevice device, WGPUQueue queue, - WGPUTextureView uTexSrc, - WGPUTextureView uTexMsk, - WGPUTextureView uTexDst, - WGPUTextureView uTexTrg); - 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 -//************************************************************************ - -class WgBindGroupOpacityPool -{ +class WgBindGroupLayouts { private: - WgBindGroupOpacity* mPool[256]{}; + WGPUDevice device{}; +public: + WGPUBindGroupLayout layoutTexSampled{}; + WGPUBindGroupLayout layoutTexScreen1WO{}; + WGPUBindGroupLayout layoutTexStrorage1WO{}; + WGPUBindGroupLayout layoutTexStrorage1RO{}; + WGPUBindGroupLayout layoutTexStrorage2RO{}; + WGPUBindGroupLayout layoutTexStrorage3RO{}; + WGPUBindGroupLayout layoutBuffer1Un{}; + WGPUBindGroupLayout layoutBuffer2Un{}; + WGPUBindGroupLayout layoutBuffer3Un{}; +public: + WGPUBindGroup createBindGroupTexSampled(WGPUSampler sampler, WGPUTextureView texView); + WGPUBindGroup createBindGroupScreen1WO(WGPUTextureView texView); + WGPUBindGroup createBindGroupStrorage1WO(WGPUTextureView texView); + WGPUBindGroup createBindGroupStrorage1RO(WGPUTextureView texView); + WGPUBindGroup createBindGroupStrorage2RO(WGPUTextureView texView0, WGPUTextureView texView1); + WGPUBindGroup createBindGroupStrorage3RO(WGPUTextureView texView0, WGPUTextureView texView1, WGPUTextureView texView2); + WGPUBindGroup createBindGroupBuffer1Un(WGPUBuffer buff); + WGPUBindGroup createBindGroupBuffer2Un(WGPUBuffer buff0, WGPUBuffer buff1); + WGPUBindGroup createBindGroupBuffer3Un(WGPUBuffer buff0, WGPUBuffer buff1, WGPUBuffer buff2); + void releaseBindGroup(WGPUBindGroup& bindGroup); 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 old mode 100644 new mode 100755 index d84cfa45..ad43f54c --- a/src/renderer/wg_engine/tvgWgCommon.cpp +++ b/src/renderer/wg_engine/tvgWgCommon.cpp @@ -20,16 +20,13 @@ * SOFTWARE. */ -#include "tvgWgCommon.h" #ifdef __EMSCRIPTEN__ #include #endif +#include "tvgWgCommon.h" +#include "tvgArray.h" #include -//***************************************************************************** -// context -//***************************************************************************** - void WgContext::initialize(WGPUInstance instance, WGPUSurface surface) { assert(instance); @@ -39,198 +36,119 @@ void WgContext::initialize(WGPUInstance instance, WGPUSurface surface) this->instance = instance; this->surface = surface; - // request adapter options - WGPURequestAdapterOptions requestAdapterOptions{}; - requestAdapterOptions.nextInChain = nullptr; - requestAdapterOptions.compatibleSurface = surface; - requestAdapterOptions.powerPreference = WGPUPowerPreference_HighPerformance; - requestAdapterOptions.forceFallbackAdapter = false; - // on adapter request ended function - auto onAdapterRequestEnded = [](WGPURequestAdapterStatus status, WGPUAdapter adapter, char const * message, void * pUserData) { - if (status != WGPURequestAdapterStatus_Success) - TVGERR("WG_RENDERER", "Adapter request: %s", message); - *((WGPUAdapter*)pUserData) = adapter; - }; // request adapter + const WGPURequestAdapterOptions requestAdapterOptions { .nextInChain = nullptr, .compatibleSurface = surface, .powerPreference = WGPUPowerPreference_HighPerformance, .forceFallbackAdapter = false }; + auto onAdapterRequestEnded = [](WGPURequestAdapterStatus status, WGPUAdapter adapter, char const * message, void * pUserData) { *((WGPUAdapter*)pUserData) = adapter; }; wgpuInstanceRequestAdapter(instance, &requestAdapterOptions, onAdapterRequestEnded, &adapter); #ifdef __EMSCRIPTEN__ while (!adapter) emscripten_sleep(10); #endif assert(adapter); - // get feature names + // get adapter and surface properties + WGPUFeatureName featureNames[32]{}; size_t featuresCount = wgpuAdapterEnumerateFeatures(adapter, featureNames); + preferredFormat = wgpuSurfaceGetPreferredFormat(surface, adapter); // request device - WGPUDeviceDescriptor deviceDesc{}; - deviceDesc.nextInChain = nullptr; - deviceDesc.label = "The device"; - deviceDesc.requiredFeatureCount = featuresCount; - deviceDesc.requiredFeatures = featureNames; - deviceDesc.requiredLimits = nullptr; - deviceDesc.defaultQueue.nextInChain = nullptr; - deviceDesc.defaultQueue.label = "The default queue"; - deviceDesc.deviceLostCallback = nullptr; - deviceDesc.deviceLostUserdata = nullptr; - // on device request ended function - auto onDeviceRequestEnded = [](WGPURequestDeviceStatus status, WGPUDevice device, char const * message, void * pUserData) { - if (status != WGPURequestDeviceStatus_Success) - TVGERR("WG_RENDERER", "Device request: %s", message); - *((WGPUDevice*)pUserData) = device; - }; - // request device + const WGPUDeviceDescriptor deviceDesc { .nextInChain = nullptr, .label = "The device", .requiredFeatureCount = featuresCount, .requiredFeatures = featureNames }; + auto onDeviceRequestEnded = [](WGPURequestDeviceStatus status, WGPUDevice device, char const * message, void * pUserData) { *((WGPUDevice*)pUserData) = device; }; wgpuAdapterRequestDevice(adapter, &deviceDesc, onDeviceRequestEnded, &device); #ifdef __EMSCRIPTEN__ while (!device) emscripten_sleep(10); #endif assert(device); - // on device error function - auto onDeviceError = [](WGPUErrorType type, char const* message, void* pUserData) { - TVGERR("WG_RENDERER", "Uncaptured device error: %s", message); - // TODO: remove direct error message - std::cout << message << std::endl; - }; - // set device error handling + // device uncaptured error callback + auto onDeviceError = [](WGPUErrorType type, char const* message, void* pUserData) { std::cout << message << std::endl; }; wgpuDeviceSetUncapturedErrorCallback(device, onDeviceError, nullptr); + // get current queue queue = wgpuDeviceGetQueue(device); assert(queue); - - // create default nearest and linear samplers + + // create shared webgpu assets + allocateBufferIndexFan(32768); samplerNearest = createSampler(WGPUFilterMode_Nearest, WGPUMipmapFilterMode_Nearest); - assert(samplerNearest); samplerLinear = createSampler(WGPUFilterMode_Linear, WGPUMipmapFilterMode_Linear); + assert(samplerNearest); assert(samplerLinear); - allocateIndexBufferFan(1024); - assert(indexBufferFan); } void WgContext::release() { - releaseBuffer(indexBufferFan); releaseSampler(samplerNearest); releaseSampler(samplerLinear); - if (device) { - wgpuDeviceDestroy(device); - wgpuDeviceRelease(device); - device = nullptr; - } - if (adapter) { - wgpuAdapterRelease(adapter); - adapter = nullptr; - } - if (surface) { - wgpuSurfaceUnconfigure(surface); - wgpuSurfaceRelease(surface); - surface = nullptr; - } - if (instance) { - wgpuInstanceRelease(instance); - instance = nullptr; - } + releaseBuffer(bufferIndexFan); + wgpuDeviceRelease(device); + wgpuAdapterRelease(adapter); } -void WgContext::executeCommandEncoder(WGPUCommandEncoder commandEncoder) +WGPUSampler WgContext::createSampler(WGPUFilterMode filter, WGPUMipmapFilterMode mipmapFilter) { - // command buffer descriptor - WGPUCommandBufferDescriptor commandBufferDesc{}; - commandBufferDesc.nextInChain = nullptr; - commandBufferDesc.label = "The command buffer"; - WGPUCommandBuffer commandsBuffer = nullptr; - commandsBuffer = wgpuCommandEncoderFinish(commandEncoder, &commandBufferDesc); - wgpuQueueSubmit(queue, 1, &commandsBuffer); - wgpuCommandBufferRelease(commandsBuffer); -} - - -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; + const WGPUSamplerDescriptor samplerDesc { + .magFilter = filter, .minFilter = filter, .mipmapFilter = mipmapFilter, + .lodMinClamp = 0.0f, .lodMaxClamp = 32.0f, .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; - 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; +WGPUTexture WgContext::createTexture(uint32_t width, uint32_t height, WGPUTextureFormat format) +{ + const WGPUTextureDescriptor textureDesc { + .usage = WGPUTextureUsage_CopyDst | WGPUTextureUsage_TextureBinding, + .dimension = WGPUTextureDimension_2D, .size = { width, height, 1 }, + .format = format, .mipLevelCount = 1, .sampleCount = 1 + }; return wgpuDeviceCreateTexture(device, &textureDesc); } -WGPUTexture WgContext::createTexture2dMS(WGPUTextureUsageFlags usage, WGPUTextureFormat format, uint32_t width, uint32_t height, uint32_t sc, char const * label) +WGPUTexture WgContext::createTexStorage(uint32_t width, uint32_t height, WGPUTextureFormat format, uint32_t sc) { - 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 = sc; - textureDesc.viewFormatCount = 0; - textureDesc.viewFormats = nullptr; + const WGPUTextureDescriptor textureDesc { + .usage = WGPUTextureUsage_CopySrc | WGPUTextureUsage_CopyDst | WGPUTextureUsage_TextureBinding | WGPUTextureUsage_StorageBinding | WGPUTextureUsage_RenderAttachment, + .dimension = WGPUTextureDimension_2D, .size = { width, height, 1 }, + .format = format, .mipLevelCount = 1, .sampleCount = sc + }; return wgpuDeviceCreateTexture(device, &textureDesc); } -WGPUTextureView WgContext::createTextureView2d(WGPUTexture texture, char const * label) +WGPUTexture WgContext::createTexStencil(uint32_t width, uint32_t height, WGPUTextureFormat format, uint32_t sc) { - 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); -}; - - -WGPUBuffer WgContext::createBuffer(WGPUBufferUsageFlags usage, uint64_t size,char const * label) -{ - WGPUBufferDescriptor bufferDesc{}; - bufferDesc.nextInChain = nullptr; - bufferDesc.label = label; - bufferDesc.usage = usage; - bufferDesc.size = size; - bufferDesc.mappedAtCreation = false; - return wgpuDeviceCreateBuffer(device, &bufferDesc); + const WGPUTextureDescriptor textureDesc { + .usage = WGPUTextureUsage_RenderAttachment, + .dimension = WGPUTextureDimension_2D, .size = { width, height, 1 }, + .format = format, .mipLevelCount = 1, .sampleCount = sc + }; + return wgpuDeviceCreateTexture(device, &textureDesc); } -void WgContext::releaseSampler(WGPUSampler& sampler) +WGPUTextureView WgContext::createTextureView(WGPUTexture texture) { - if (sampler) { - wgpuSamplerRelease(sampler); - sampler = nullptr; + const WGPUTextureViewDescriptor textureViewDesc { + .format = wgpuTextureGetFormat(texture), + .dimension = WGPUTextureViewDimension_2D, + .baseMipLevel = 0, + .mipLevelCount = 1, + .baseArrayLayer = 0, + .arrayLayerCount = 1, + .aspect = WGPUTextureAspect_All + }; + return wgpuTextureCreateView(texture, &textureViewDesc); +} + + +void WgContext::releaseTextureView(WGPUTextureView& textureView) +{ + if (textureView) { + wgpuTextureViewRelease(textureView); + textureView = nullptr; } } @@ -246,15 +164,89 @@ void WgContext::releaseTexture(WGPUTexture& texture) } -void WgContext::releaseTextureView(WGPUTextureView& textureView) +void WgContext::releaseSampler(WGPUSampler& sampler) { - if (textureView) { - wgpuTextureViewRelease(textureView); - textureView = nullptr; + if (sampler) { + wgpuSamplerRelease(sampler); + sampler = nullptr; } } +bool WgContext::allocateBufferUniform(WGPUBuffer& buffer, const void* data, uint64_t size) +{ + if ((buffer) && (wgpuBufferGetSize(buffer) >= size)) + wgpuQueueWriteBuffer(queue, buffer, 0, data, size); + else { + releaseBuffer(buffer); + const WGPUBufferDescriptor bufferDesc { .usage = WGPUBufferUsage_CopyDst | WGPUBufferUsage_Uniform, .size = size }; + buffer = wgpuDeviceCreateBuffer(device, &bufferDesc); + wgpuQueueWriteBuffer(queue, buffer, 0, data, size); + return true; + } + return false; +} + + +bool WgContext::allocateBufferVertex(WGPUBuffer& buffer, const float* data, uint64_t size) +{ + if ((buffer) && (wgpuBufferGetSize(buffer) >= size)) + wgpuQueueWriteBuffer(queue, buffer, 0, data, size); + else { + releaseBuffer(buffer); + const WGPUBufferDescriptor bufferDesc { + .usage = WGPUBufferUsage_CopyDst | WGPUBufferUsage_Vertex, + .size = size > WG_VERTEX_BUFFER_MIN_SIZE ? size : WG_VERTEX_BUFFER_MIN_SIZE + }; + buffer = wgpuDeviceCreateBuffer(device, &bufferDesc); + wgpuQueueWriteBuffer(queue, buffer, 0, data, size); + return true; + } + return false; +} + + +bool WgContext::allocateBufferIndex(WGPUBuffer& buffer, const uint32_t* data, uint64_t size) +{ + if ((buffer) && (wgpuBufferGetSize(buffer) >= size)) + wgpuQueueWriteBuffer(queue, buffer, 0, data, size); + else { + releaseBuffer(buffer); + const WGPUBufferDescriptor bufferDesc { + .usage = WGPUBufferUsage_CopyDst | WGPUBufferUsage_Index, + .size = size > WG_INDEX_BUFFER_MIN_SIZE ? size : WG_INDEX_BUFFER_MIN_SIZE + }; + buffer = wgpuDeviceCreateBuffer(device, &bufferDesc); + wgpuQueueWriteBuffer(queue, buffer, 0, data, size); + return true; + } + return false; +} + + +bool WgContext::allocateBufferIndexFan(uint64_t vertexCount) +{ + uint64_t indexCount = (vertexCount - 2) * 3; + if ((!bufferIndexFan) || (wgpuBufferGetSize(bufferIndexFan) < indexCount * sizeof(uint32_t))) { + tvg::Array indexes(indexCount); + for (size_t i = 0; i < vertexCount - 2; i++) { + indexes.push(0); + indexes.push(i + 1); + indexes.push(i + 2); + } + releaseBuffer(bufferIndexFan); + WGPUBufferDescriptor bufferDesc{ + .usage = WGPUBufferUsage_CopyDst | WGPUBufferUsage_Index, + .size = indexCount * sizeof(uint32_t) + }; + bufferIndexFan = wgpuDeviceCreateBuffer(device, &bufferDesc); + wgpuQueueWriteBuffer(queue, bufferIndexFan, 0, &indexes[0], indexCount * sizeof(uint32_t)); + return true; + } + return false; +} + + void WgContext::releaseBuffer(WGPUBuffer& buffer) { if (buffer) { @@ -263,540 +255,3 @@ void WgContext::releaseBuffer(WGPUBuffer& buffer) buffer = nullptr; } } - - -void WgContext::allocateVertexBuffer(WGPUBuffer& buffer, const void *data, uint64_t size) -{ - if ((buffer) && (wgpuBufferGetSize(buffer) >= size)) - wgpuQueueWriteBuffer(queue, buffer, 0, data, size); - else { - // create new buffer and upload data - releaseBuffer(buffer); - WGPUBufferDescriptor bufferDesc{}; - bufferDesc.nextInChain = nullptr; - bufferDesc.label = "The vertex buffer"; - bufferDesc.usage = WGPUBufferUsage_CopyDst | WGPUBufferUsage_Vertex; - bufferDesc.size = size > WG_VERTEX_BUFFER_MIN_SIZE ? size : WG_VERTEX_BUFFER_MIN_SIZE; - bufferDesc.mappedAtCreation = false; - buffer = wgpuDeviceCreateBuffer(device, &bufferDesc); - wgpuQueueWriteBuffer(queue, buffer, 0, data, size); - } -} - - -void WgContext::allocateIndexBuffer(WGPUBuffer& buffer, const void *data, uint64_t size) -{ - if ((buffer) && (wgpuBufferGetSize(buffer) >= size)) - wgpuQueueWriteBuffer(queue, buffer, 0, data, size); - else { - // create new buffer and upload data - releaseBuffer(buffer); - WGPUBufferDescriptor bufferDesc{}; - bufferDesc.nextInChain = nullptr; - bufferDesc.label = "The index buffer"; - bufferDesc.usage = WGPUBufferUsage_CopyDst | WGPUBufferUsage_Index; - bufferDesc.size = size > WG_INDEX_BUFFER_MIN_SIZE ? size : WG_INDEX_BUFFER_MIN_SIZE; - bufferDesc.mappedAtCreation = false; - buffer = wgpuDeviceCreateBuffer(device, &bufferDesc); - wgpuQueueWriteBuffer(queue, buffer, 0, data, size); - } -} - - -void WgContext::allocateIndexBufferFan(uint64_t vertsCount) -{ - assert(vertsCount >= 3); - uint64_t indexCount = (vertsCount - 2) * 3; - if ((!indexBufferFan) || (wgpuBufferGetSize(indexBufferFan) < indexCount * sizeof(uint32_t))) { - Array indexes(indexCount); - for (size_t i = 0; i < vertsCount - 2; i++) { - indexes.push(0); - indexes.push(i + 1); - indexes.push(i + 2); - } - releaseBuffer(indexBufferFan); - WGPUBufferDescriptor bufferDesc{}; - bufferDesc.nextInChain = nullptr; - bufferDesc.label = "The index buffer"; - bufferDesc.usage = WGPUBufferUsage_CopyDst | WGPUBufferUsage_Index; - bufferDesc.size = indexCount * sizeof(uint32_t); - bufferDesc.mappedAtCreation = false; - indexBufferFan = wgpuDeviceCreateBuffer(device, &bufferDesc); - wgpuQueueWriteBuffer(queue, indexBufferFan, 0, &indexes[0], indexCount * sizeof(uint32_t)); - } -} - - -void WgContext::releaseVertexBuffer(WGPUBuffer& buffer) -{ - releaseBuffer(buffer); -} - - -void WgContext::releaseIndexBuffer(WGPUBuffer& buffer) -{ - releaseBuffer(buffer); -} - -//***************************************************************************** -// bind group -//***************************************************************************** - -void WgBindGroup::set(WGPURenderPassEncoder encoder, uint32_t groupIndex) -{ - wgpuRenderPassEncoderSetBindGroup(encoder, groupIndex, mBindGroup, 0, nullptr); -} - - -void WgBindGroup::set(WGPUComputePassEncoder encoder, uint32_t groupIndex) -{ - wgpuComputePassEncoderSetBindGroup(encoder, groupIndex, mBindGroup, 0, nullptr); -} - - -WGPUBindGroupEntry WgBindGroup::makeBindGroupEntryBuffer(uint32_t binding, WGPUBuffer buffer) -{ - WGPUBindGroupEntry bindGroupEntry{}; - bindGroupEntry.nextInChain = nullptr; - bindGroupEntry.binding = binding; - bindGroupEntry.buffer = buffer; - bindGroupEntry.offset = 0; - bindGroupEntry.size = wgpuBufferGetSize(buffer); - bindGroupEntry.sampler = nullptr; - bindGroupEntry.textureView = nullptr; - return bindGroupEntry; -} - - -WGPUBindGroupEntry WgBindGroup::makeBindGroupEntrySampler(uint32_t binding, WGPUSampler sampler) -{ - WGPUBindGroupEntry bindGroupEntry{}; - bindGroupEntry.nextInChain = nullptr; - bindGroupEntry.binding = binding; - bindGroupEntry.buffer = nullptr; - bindGroupEntry.offset = 0; - bindGroupEntry.size = 0; - bindGroupEntry.sampler = sampler; - bindGroupEntry.textureView = nullptr; - return bindGroupEntry; -} - - -WGPUBindGroupEntry WgBindGroup::makeBindGroupEntryTextureView(uint32_t binding, WGPUTextureView textureView) -{ - WGPUBindGroupEntry bindGroupEntry{}; - bindGroupEntry.nextInChain = nullptr; - bindGroupEntry.binding = binding; - bindGroupEntry.buffer = nullptr; - bindGroupEntry.offset = 0; - bindGroupEntry.size = 0; - bindGroupEntry.sampler = nullptr; - bindGroupEntry.textureView = textureView; - return bindGroupEntry; -} - - -WGPUBindGroupLayoutEntry WgBindGroup::makeBindGroupLayoutEntryBuffer(uint32_t binding) -{ - WGPUBindGroupLayoutEntry bindGroupLayoutEntry{}; - bindGroupLayoutEntry.nextInChain = nullptr; - bindGroupLayoutEntry.binding = binding; - bindGroupLayoutEntry.visibility = WGPUShaderStage_Vertex | WGPUShaderStage_Fragment | WGPUShaderStage_Compute; - bindGroupLayoutEntry.buffer.nextInChain = nullptr; - bindGroupLayoutEntry.buffer.type = WGPUBufferBindingType_Uniform; - bindGroupLayoutEntry.buffer.hasDynamicOffset = false; - bindGroupLayoutEntry.buffer.minBindingSize = 0; - return bindGroupLayoutEntry; -} - - -WGPUBindGroupLayoutEntry WgBindGroup::makeBindGroupLayoutEntrySampler(uint32_t binding) -{ - WGPUBindGroupLayoutEntry bindGroupLayoutEntry{}; - bindGroupLayoutEntry.nextInChain = nullptr; - bindGroupLayoutEntry.binding = binding; - bindGroupLayoutEntry.visibility = WGPUShaderStage_Fragment; - bindGroupLayoutEntry.sampler.nextInChain = nullptr; - bindGroupLayoutEntry.sampler.type = WGPUSamplerBindingType_Filtering; - return bindGroupLayoutEntry; -} - - -WGPUBindGroupLayoutEntry WgBindGroup::makeBindGroupLayoutEntryTexture(uint32_t binding) -{ - WGPUBindGroupLayoutEntry bindGroupLayoutEntry{}; - bindGroupLayoutEntry.nextInChain = nullptr; - bindGroupLayoutEntry.binding = binding; - bindGroupLayoutEntry.visibility = WGPUShaderStage_Fragment | WGPUShaderStage_Compute; - bindGroupLayoutEntry.texture.nextInChain = nullptr; - bindGroupLayoutEntry.texture.sampleType = WGPUTextureSampleType_Float; - bindGroupLayoutEntry.texture.viewDimension = WGPUTextureViewDimension_2D; - bindGroupLayoutEntry.texture.multisampled = false; - return bindGroupLayoutEntry; -} - - -WGPUBindGroupLayoutEntry WgBindGroup::makeBindGroupLayoutEntryStorage(uint32_t binding, WGPUStorageTextureAccess access, WGPUTextureFormat format) -{ - WGPUBindGroupLayoutEntry bindGroupLayoutEntry{}; - bindGroupLayoutEntry.nextInChain = nullptr; - bindGroupLayoutEntry.binding = binding; - bindGroupLayoutEntry.visibility = WGPUShaderStage_Fragment | WGPUShaderStage_Compute; - bindGroupLayoutEntry.storageTexture.nextInChain = nullptr; - bindGroupLayoutEntry.storageTexture.access = access; - bindGroupLayoutEntry.storageTexture.format = format; - bindGroupLayoutEntry.storageTexture.viewDimension = WGPUTextureViewDimension_2D; - return bindGroupLayoutEntry; -} - - -WGPUBuffer WgBindGroup::createBuffer(WGPUDevice device, WGPUQueue queue, const void *data, size_t size) -{ - WGPUBufferDescriptor bufferDescriptor{}; - bufferDescriptor.nextInChain = nullptr; - bufferDescriptor.label = "The uniform buffer"; - bufferDescriptor.usage = WGPUBufferUsage_CopyDst | WGPUBufferUsage_Uniform; - bufferDescriptor.size = size; - bufferDescriptor.mappedAtCreation = false; - WGPUBuffer buffer = wgpuDeviceCreateBuffer(device, &bufferDescriptor); - assert(buffer); - wgpuQueueWriteBuffer(queue, buffer, 0, data, size); - return buffer; -} - - -WGPUBindGroup WgBindGroup::createBindGroup(WGPUDevice device, WGPUBindGroupLayout layout, const WGPUBindGroupEntry* bindGroupEntries, uint32_t count) -{ - WGPUBindGroupDescriptor bindGroupDesc{}; - bindGroupDesc.nextInChain = nullptr; - bindGroupDesc.label = "The binding group"; - bindGroupDesc.layout = layout; - bindGroupDesc.entryCount = count; - bindGroupDesc.entries = bindGroupEntries; - return wgpuDeviceCreateBindGroup(device, &bindGroupDesc); -} - - -WGPUBindGroupLayout WgBindGroup::createBindGroupLayout(WGPUDevice device, const WGPUBindGroupLayoutEntry* bindGroupLayoutEntries, uint32_t count) -{ - WGPUBindGroupLayoutDescriptor bindGroupLayoutDesc{}; - bindGroupLayoutDesc.nextInChain = nullptr; - bindGroupLayoutDesc.label = "The bind group layout"; - bindGroupLayoutDesc.entryCount = count; - bindGroupLayoutDesc.entries = bindGroupLayoutEntries; // @binding - return wgpuDeviceCreateBindGroupLayout(device, &bindGroupLayoutDesc); -} - - -void WgBindGroup::releaseBuffer(WGPUBuffer& buffer) -{ - if (buffer) { - wgpuBufferDestroy(buffer); - wgpuBufferRelease(buffer); - buffer = nullptr; - } -} - - -void WgBindGroup::releaseBindGroup(WGPUBindGroup& bindGroup) -{ - if (bindGroup) wgpuBindGroupRelease(bindGroup); - bindGroup = nullptr; -} - - -void WgBindGroup::releaseBindGroupLayout(WGPUBindGroupLayout& bindGroupLayout) -{ - if (bindGroupLayout) wgpuBindGroupLayoutRelease(bindGroupLayout); - bindGroupLayout = nullptr; -} - -//***************************************************************************** -// pipeline -//***************************************************************************** - -void WgPipeline::release() -{ - destroyShaderModule(mShaderModule); - destroyPipelineLayout(mPipelineLayout); -} - - -WGPUPipelineLayout WgPipeline::createPipelineLayout(WGPUDevice device, const WGPUBindGroupLayout* bindGroupLayouts, uint32_t count) -{ - WGPUPipelineLayoutDescriptor pipelineLayoutDesc{}; - pipelineLayoutDesc.nextInChain = nullptr; - pipelineLayoutDesc.label = "The pipeline layout"; - pipelineLayoutDesc.bindGroupLayoutCount = count; - pipelineLayoutDesc.bindGroupLayouts = bindGroupLayouts; - return wgpuDeviceCreatePipelineLayout(device, &pipelineLayoutDesc); -} - - -WGPUShaderModule WgPipeline::createShaderModule(WGPUDevice device, const char* code, const char* label) -{ - WGPUShaderModuleWGSLDescriptor shaderModuleWGSLDesc{}; - shaderModuleWGSLDesc.chain.next = nullptr; - shaderModuleWGSLDesc.chain.sType = WGPUSType_ShaderModuleWGSLDescriptor; - shaderModuleWGSLDesc.code = code; - WGPUShaderModuleDescriptor shaderModuleDesc{}; - shaderModuleDesc.nextInChain = &shaderModuleWGSLDesc.chain; - shaderModuleDesc.label = label; - #ifndef __EMSCRIPTEN__ - shaderModuleDesc.hintCount = 0; - shaderModuleDesc.hints = nullptr; - #endif - return wgpuDeviceCreateShaderModule(device, &shaderModuleDesc); -} - - -void WgPipeline::destroyPipelineLayout(WGPUPipelineLayout& pipelineLayout) -{ - if (pipelineLayout) wgpuPipelineLayoutRelease(pipelineLayout); - pipelineLayout = nullptr; -} - - -void WgPipeline::destroyShaderModule(WGPUShaderModule& shaderModule) -{ - if (shaderModule) wgpuShaderModuleRelease(shaderModule); - shaderModule = nullptr; -} - -//***************************************************************************** -// render pipeline -//***************************************************************************** - -void WgRenderPipeline::allocate(WGPUDevice device, WgPipelineBlendType blendType, WGPUColorWriteMask writeMask, - WGPUVertexBufferLayout vertexBufferLayouts[], uint32_t attribsCount, - WGPUBindGroupLayout bindGroupLayouts[], uint32_t bindGroupsCount, - WGPUCompareFunction compareFront, WGPUStencilOperation operationFront, - WGPUCompareFunction compareBack, WGPUStencilOperation operationBack, - const char* shaderSource, const char* shaderLabel, const char* pipelineLabel) -{ - mShaderModule = createShaderModule(device, shaderSource, shaderLabel); - assert(mShaderModule); - - mPipelineLayout = createPipelineLayout(device, bindGroupLayouts, bindGroupsCount); - assert(mPipelineLayout); - - mRenderPipeline = createRenderPipeline(device, blendType, writeMask, - vertexBufferLayouts, attribsCount, - compareFront, operationFront, - compareBack, operationBack, - mPipelineLayout, mShaderModule, pipelineLabel); - assert(mRenderPipeline); -} - - -void WgRenderPipeline::release() -{ - destroyRenderPipeline(mRenderPipeline); - WgPipeline::release(); -} - - -void WgRenderPipeline::set(WGPURenderPassEncoder renderPassEncoder) -{ - wgpuRenderPassEncoderSetPipeline(renderPassEncoder, mRenderPipeline); -}; - - -WGPUBlendState WgRenderPipeline::makeBlendState(WgPipelineBlendType blendType) -{ - WGPUBlendState blendState{}; - // srcover - if (blendType == WgPipelineBlendType::SrcOver) { - blendState.color.operation = WGPUBlendOperation_Add; - blendState.color.srcFactor = WGPUBlendFactor_One; - blendState.color.dstFactor = WGPUBlendFactor_Zero; - // normal - } else if (blendType == WgPipelineBlendType::Normal) { - blendState.color.operation = WGPUBlendOperation_Add; - blendState.color.srcFactor = WGPUBlendFactor_One; - blendState.color.dstFactor = WGPUBlendFactor_OneMinusSrcAlpha; - // custom - } else if (blendType == WgPipelineBlendType::Custom) { - blendState.color.operation = WGPUBlendOperation_Add; - blendState.color.srcFactor = WGPUBlendFactor_One; - blendState.color.dstFactor = WGPUBlendFactor_Zero; - } - blendState.alpha = blendState.color; - return blendState; -} - - -WGPUColorTargetState WgRenderPipeline::makeColorTargetState(const WGPUBlendState* blendState, const WGPUColorWriteMask writeMask) -{ - WGPUColorTargetState colorTargetState{}; - colorTargetState.nextInChain = nullptr; - colorTargetState.format = WGPUTextureFormat_RGBA8Unorm; - colorTargetState.blend = blendState; - colorTargetState.writeMask = writeMask; - return colorTargetState; -} - - -WGPUVertexBufferLayout WgRenderPipeline::makeVertexBufferLayout(const WGPUVertexAttribute* vertexAttributes, uint32_t count, uint64_t stride) -{ - WGPUVertexBufferLayout vertexBufferLayoutPos{}; - vertexBufferLayoutPos.arrayStride = stride; - vertexBufferLayoutPos.stepMode = WGPUVertexStepMode_Vertex; - vertexBufferLayoutPos.attributeCount = count; - vertexBufferLayoutPos.attributes = vertexAttributes; - return vertexBufferLayoutPos; -} - - -WGPUVertexState WgRenderPipeline::makeVertexState(WGPUShaderModule shaderModule, const WGPUVertexBufferLayout* buffers, uint32_t count) -{ - WGPUVertexState vertexState{}; - vertexState.nextInChain = nullptr; - vertexState.module = shaderModule; - vertexState.entryPoint = "vs_main"; - vertexState.constantCount = 0; - vertexState.constants = nullptr; - vertexState.bufferCount = count; - vertexState.buffers = buffers; - return vertexState; -} - - -WGPUPrimitiveState WgRenderPipeline::makePrimitiveState() -{ - WGPUPrimitiveState primitiveState{}; - primitiveState.nextInChain = nullptr; - primitiveState.topology = WGPUPrimitiveTopology_TriangleList; - primitiveState.stripIndexFormat = WGPUIndexFormat_Undefined; - primitiveState.frontFace = WGPUFrontFace_CCW; - primitiveState.cullMode = WGPUCullMode_None; - return primitiveState; -} - - -WGPUDepthStencilState WgRenderPipeline::makeDepthStencilState(WGPUCompareFunction compareFront, WGPUStencilOperation operationFront, WGPUCompareFunction compareBack, WGPUStencilOperation operationBack) -{ - WGPUDepthStencilState depthStencilState{}; - depthStencilState.nextInChain = nullptr; - depthStencilState.format = WGPUTextureFormat_Stencil8; - depthStencilState.depthWriteEnabled = false; - depthStencilState.depthCompare = WGPUCompareFunction_Always; - depthStencilState.stencilFront.compare = compareFront; - depthStencilState.stencilFront.failOp = operationFront; - depthStencilState.stencilFront.depthFailOp = operationFront; - depthStencilState.stencilFront.passOp = operationFront; - depthStencilState.stencilBack.compare = compareBack; - depthStencilState.stencilBack.failOp = operationBack; - depthStencilState.stencilBack.depthFailOp = operationBack; - depthStencilState.stencilBack.passOp = operationBack; - depthStencilState.stencilReadMask = 0xFFFFFFFF; - depthStencilState.stencilWriteMask = 0xFFFFFFFF; - depthStencilState.depthBias = 0; - depthStencilState.depthBiasSlopeScale = 0.0f; - depthStencilState.depthBiasClamp = 0.0f; - return depthStencilState; -} - - -WGPUMultisampleState WgRenderPipeline::makeMultisampleState() -{ - WGPUMultisampleState multisampleState{}; - multisampleState.nextInChain = nullptr; - multisampleState.count = 1; - multisampleState.mask = 0xFFFFFFFF; - multisampleState.alphaToCoverageEnabled = false; - return multisampleState; -} - - -WGPUFragmentState WgRenderPipeline::makeFragmentState(WGPUShaderModule shaderModule, WGPUColorTargetState* targets, uint32_t size) -{ - WGPUFragmentState fragmentState{}; - fragmentState.nextInChain = nullptr; - fragmentState.module = shaderModule; - fragmentState.entryPoint = "fs_main"; - fragmentState.constantCount = 0; - fragmentState.constants = nullptr; - fragmentState.targetCount = size; - fragmentState.targets = targets; - return fragmentState; -} - - -WGPURenderPipeline WgRenderPipeline::createRenderPipeline(WGPUDevice device, WgPipelineBlendType blendType, WGPUColorWriteMask writeMask, - WGPUVertexBufferLayout vertexBufferLayouts[], uint32_t attribsCount, - WGPUCompareFunction compareFront, WGPUStencilOperation operationFront, - WGPUCompareFunction compareBack, WGPUStencilOperation operationBack, - WGPUPipelineLayout pipelineLayout, WGPUShaderModule shaderModule, - const char* pipelineName) -{ - WGPUBlendState blendState = makeBlendState(blendType); - WGPUColorTargetState colorTargetStates[] = { - makeColorTargetState(&blendState, writeMask) - }; - - WGPUVertexState vertexState = makeVertexState(shaderModule, vertexBufferLayouts, attribsCount); - WGPUPrimitiveState primitiveState = makePrimitiveState(); - WGPUDepthStencilState depthStencilState = makeDepthStencilState(compareFront, operationFront, compareBack, operationBack); - WGPUMultisampleState multisampleState = makeMultisampleState(); - WGPUFragmentState fragmentState = makeFragmentState(shaderModule, colorTargetStates, 1); - - WGPURenderPipelineDescriptor renderPipelineDesc{}; - renderPipelineDesc.nextInChain = nullptr; - renderPipelineDesc.label = pipelineName; - renderPipelineDesc.layout = pipelineLayout; - renderPipelineDesc.vertex = vertexState; - renderPipelineDesc.primitive = primitiveState; - renderPipelineDesc.depthStencil = &depthStencilState; - renderPipelineDesc.multisample = multisampleState; - renderPipelineDesc.fragment = &fragmentState; - return wgpuDeviceCreateRenderPipeline(device, &renderPipelineDesc); -} - - -void WgRenderPipeline::destroyRenderPipeline(WGPURenderPipeline& renderPipeline) -{ - if (renderPipeline) wgpuRenderPipelineRelease(renderPipeline); - renderPipeline = nullptr; -} - -//***************************************************************************** -// 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 old mode 100644 new mode 100755 index d0cdef3a..b6095caa --- a/src/renderer/wg_engine/tvgWgCommon.h +++ b/src/renderer/wg_engine/tvgWgCommon.h @@ -25,151 +25,51 @@ #include #include -#include "tvgCommon.h" -#include "tvgRender.h" #define WG_VERTEX_BUFFER_MIN_SIZE 2048 #define WG_INDEX_BUFFER_MIN_SIZE 2048 -struct WgCompositor: public Compositor { - BlendMethod blendMethod; -}; - -enum class WgPipelineBlendType { - SrcOver = 0, - Normal, - Custom -}; - -struct WgPipelines; +class WgPipelines; struct WgContext { + // external webgpu handles WGPUInstance instance{}; WGPUSurface surface{}; + // common webgpu handles WGPUAdapter adapter{}; WGPUDevice device{}; WGPUQueue queue{}; - - WGPUFeatureName featureNames[32]{}; - WGPUAdapterProperties adapterProperties{}; - WGPUSupportedLimits supportedLimits{}; - + WGPUTextureFormat preferredFormat{}; + // external handles (do not release) + WgPipelines* pipelines{}; + // shared webgpu assets + WGPUBuffer bufferIndexFan{}; WGPUSampler samplerNearest{}; WGPUSampler samplerLinear{}; - WGPUBuffer indexBufferFan{}; - WgPipelines* pipelines{}; // external handle (do not release) - void initialize(WGPUInstance instance, WGPUSurface surface); void release(); + + // create common objects + WGPUSampler createSampler(WGPUFilterMode filter, WGPUMipmapFilterMode mipmapFilter); + WGPUTexture createTexture(uint32_t width, uint32_t height, WGPUTextureFormat format); + WGPUTexture createTexStorage(uint32_t width, uint32_t height, WGPUTextureFormat format, uint32_t sc = 1); + WGPUTexture createTexStencil(uint32_t width, uint32_t height, WGPUTextureFormat format, uint32_t sc = 1); + WGPUTextureView createTextureView(WGPUTexture texture); - void executeCommandEncoder(WGPUCommandEncoder commandEncoder); - - WGPUSampler createSampler(WGPUFilterMode minFilter, WGPUMipmapFilterMode mipmapFilter); - WGPUTexture createTexture2d(WGPUTextureUsageFlags usage, WGPUTextureFormat format, uint32_t width, uint32_t height, char const * label); - WGPUTexture createTexture2dMS(WGPUTextureUsageFlags usage, WGPUTextureFormat format, uint32_t width, uint32_t height, uint32_t sc, char const * label); - WGPUTextureView createTextureView2d(WGPUTexture texture, char const * label); - WGPUBuffer createBuffer(WGPUBufferUsageFlags usage, uint64_t size, char const * label); - - void releaseSampler(WGPUSampler& sampler); - void releaseTexture(WGPUTexture& texture); + // release common objects void releaseTextureView(WGPUTextureView& textureView); + void releaseTexture(WGPUTexture& texture); + void releaseSampler(WGPUSampler& sampler); + + // create buffer objects (return true, if buffer handle was changed) + bool allocateBufferUniform(WGPUBuffer& buffer, const void* data, uint64_t size); + bool allocateBufferVertex(WGPUBuffer& buffer, const float* data, uint64_t size); + bool allocateBufferIndex(WGPUBuffer& buffer, const uint32_t* data, uint64_t size); + bool allocateBufferIndexFan(uint64_t vertexCount); + + // release buffer objects void releaseBuffer(WGPUBuffer& buffer); - - void allocateVertexBuffer(WGPUBuffer& buffer, const void *data, uint64_t size); - void allocateIndexBuffer(WGPUBuffer& buffer, const void *data, uint64_t size); - void allocateIndexBufferFan(uint64_t size); - void releaseVertexBuffer(WGPUBuffer& buffer); - void releaseIndexBuffer(WGPUBuffer& buffer); -}; - -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); - static WGPUBindGroupEntry makeBindGroupEntryTextureView(uint32_t binding, WGPUTextureView textureView); - - static WGPUBindGroupLayoutEntry makeBindGroupLayoutEntryBuffer(uint32_t binding); - static WGPUBindGroupLayoutEntry makeBindGroupLayoutEntrySampler(uint32_t binding); - static WGPUBindGroupLayoutEntry makeBindGroupLayoutEntryTexture(uint32_t binding); - static WGPUBindGroupLayoutEntry makeBindGroupLayoutEntryStorage(uint32_t binding, WGPUStorageTextureAccess access, WGPUTextureFormat format); - - 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); - static WGPUBindGroupLayout createBindGroupLayout(WGPUDevice device, const WGPUBindGroupLayoutEntry* bindGroupLayoutEntries, uint32_t count); - - static void releaseBuffer(WGPUBuffer& buffer); - static void releaseBindGroup(WGPUBindGroup& bindGroup); - static void releaseBindGroupLayout(WGPUBindGroupLayout& bindGroupLayout); -}; - -struct WgPipeline -{ -protected: - WGPUPipelineLayout mPipelineLayout{}; - WGPUShaderModule mShaderModule{}; -public: - virtual ~WgPipeline() {} - virtual void initialize(WGPUDevice device) = 0; - virtual void release(); - - static WGPUPipelineLayout createPipelineLayout(WGPUDevice device, const WGPUBindGroupLayout* bindGroupLayouts, uint32_t count); - static WGPUShaderModule createShaderModule(WGPUDevice device, const char* code, const char* label); - static void destroyPipelineLayout(WGPUPipelineLayout& pipelineLayout); - static void destroyShaderModule(WGPUShaderModule& shaderModule); -}; - -struct WgRenderPipeline: public WgPipeline -{ -protected: - WGPURenderPipeline mRenderPipeline{}; - void allocate(WGPUDevice device, WgPipelineBlendType blendType, WGPUColorWriteMask writeMask, - WGPUVertexBufferLayout vertexBufferLayouts[], uint32_t attribsCount, - WGPUBindGroupLayout bindGroupLayouts[], uint32_t bindGroupsCount, - WGPUCompareFunction compareFront, WGPUStencilOperation operationFront, - WGPUCompareFunction compareBack, WGPUStencilOperation operationBack, - const char* shaderSource, const char* shaderLabel, const char* pipelineLabel); -public: - void release() override; - void set(WGPURenderPassEncoder renderPassEncoder); - - static WGPUBlendState makeBlendState(WgPipelineBlendType blendType); - static WGPUColorTargetState makeColorTargetState(const WGPUBlendState* blendState, const WGPUColorWriteMask writeMask); - static WGPUVertexBufferLayout makeVertexBufferLayout(const WGPUVertexAttribute* vertexAttributes, uint32_t count, uint64_t stride); - static WGPUVertexState makeVertexState(WGPUShaderModule shaderModule, const WGPUVertexBufferLayout* buffers, uint32_t count); - static WGPUPrimitiveState makePrimitiveState(); - static WGPUDepthStencilState makeDepthStencilState(WGPUCompareFunction compareFront, WGPUStencilOperation operationFront, WGPUCompareFunction compareBack, WGPUStencilOperation operationBack); - static WGPUMultisampleState makeMultisampleState(); - static WGPUFragmentState makeFragmentState(WGPUShaderModule shaderModule, WGPUColorTargetState* targets, uint32_t size); - - static WGPURenderPipeline createRenderPipeline(WGPUDevice device, WgPipelineBlendType blendType, WGPUColorWriteMask writeMask, - WGPUVertexBufferLayout vertexBufferLayouts[], uint32_t attribsCount, - WGPUCompareFunction compareFront, WGPUStencilOperation operationFront, - WGPUCompareFunction compareBack, WGPUStencilOperation operationBack, - WGPUPipelineLayout pipelineLayout, WGPUShaderModule shaderModule, - const char* pipelineLabel); - static void destroyRenderPipeline(WGPURenderPipeline& renderPipeline); -}; - -#define WG_COMPUTE_WORKGROUP_SIZE_X 8 -#define WG_COMPUTE_WORKGROUP_SIZE_Y 8 - -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/tvgWgCompositor.cpp b/src/renderer/wg_engine/tvgWgCompositor.cpp new file mode 100755 index 00000000..a1d5368d --- /dev/null +++ b/src/renderer/wg_engine/tvgWgCompositor.cpp @@ -0,0 +1,509 @@ +/* + * Copyright (c) 2024 the ThorVG project. All rights reserved. + + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ + +#include "tvgWgCompositor.h" +#include "tvgWgShaderTypes.h" + +void WgCompositor::initialize(WgContext& context, uint32_t width, uint32_t height) +{ + // pipelines (external handle, do not release) + pipelines = context.pipelines; + // store render target dimensions + this->width = width; + this->height = height; + // allocate global stencil buffer handles + texStencil = context.createTexStencil(width, height, WGPUTextureFormat_Stencil8); + texViewStencil = context.createTextureView(texStencil); + // allocate global view matrix handles + WgShaderTypeMat4x4f viewMat(width, height); + context.allocateBufferUniform(bufferViewMat, &viewMat, sizeof(viewMat)); + bindGroupViewMat = pipelines->layouts.createBindGroupBuffer1Un(bufferViewMat); + // initialize opacity pool + for (uint32_t i = 0; i < 256; i++) { + float opacity = i / 255.0f; + context.allocateBufferUniform(bufferOpacities[i], &opacity, sizeof(float)); + bindGroupOpacities[i] = pipelines->layouts.createBindGroupBuffer1Un(bufferOpacities[i]); + } + // initialize intermediate render storages + storageClipPath.initialize(context, width, height); + storageInterm.initialize(context, width, height); + storageDstCopy.initialize(context, width, height); +} + + +void WgCompositor::release(WgContext& context) +{ + // release intermediate render storages + storageInterm.release(context); + storageDstCopy.release(context); + storageClipPath.release(context); + // release opacity pool + for (uint32_t i = 0; i < 256; i++) { + context.pipelines->layouts.releaseBindGroup(bindGroupOpacities[i]); + context.releaseBuffer(bufferOpacities[i]); + } + // release global view matrix handles + context.pipelines->layouts.releaseBindGroup(bindGroupViewMat); + context.releaseBuffer(bufferViewMat); + // release global stencil buffer handles + context.releaseTextureView(texViewStencil); + context.releaseTexture(texStencil); + height = 0; + width = 0; + pipelines = nullptr; +} + + +WgPipelineBlendType WgCompositor::blendMethodToBlendType(BlendMethod blendMethod) +{ + if (blendMethod == BlendMethod::Normal) return WgPipelineBlendType::Normal; + if (blendMethod == BlendMethod::SrcOver) return WgPipelineBlendType::SrcOver; + return WgPipelineBlendType::Custom; +} + + +void WgCompositor::beginRenderPass(WGPUCommandEncoder commandEncoder, WgRenderStorage* target, bool clear) +{ + assert(commandEncoder); + assert(target); + this->currentTarget = target; + this->commandEncoder = commandEncoder; + WGPURenderPassDepthStencilAttachment depthStencilAttachment{ .view = texViewStencil, .stencilLoadOp = WGPULoadOp_Clear, .stencilStoreOp = WGPUStoreOp_Discard }; + WGPURenderPassColorAttachment colorAttachment{}; + colorAttachment.view = target->texView, + colorAttachment.loadOp = clear ? WGPULoadOp_Clear : WGPULoadOp_Load, + colorAttachment.storeOp = WGPUStoreOp_Store; + #ifdef __EMSCRIPTEN__ + colorAttachment.depthSlice = WGPU_DEPTH_SLICE_UNDEFINED; + #endif + WGPURenderPassDescriptor renderPassDesc{ .colorAttachmentCount = 1, .colorAttachments = &colorAttachment, .depthStencilAttachment = &depthStencilAttachment }; + renderPassEncoder = wgpuCommandEncoderBeginRenderPass(commandEncoder, &renderPassDesc); + assert(renderPassEncoder); +} + + +void WgCompositor::endRenderPass() +{ + assert(renderPassEncoder); + wgpuRenderPassEncoderEnd(renderPassEncoder); + wgpuRenderPassEncoderRelease(renderPassEncoder); + this->renderPassEncoder = nullptr; + this->currentTarget = nullptr; +} + + +void WgCompositor::renderClipPath(WgContext& context, WgRenderDataPaint* renderData, WgRenderStorage* dst) +{ + assert(renderData); + if (renderData->clips.count == 0) return; + // store current render pass + WgRenderStorage *target = currentTarget; + endRenderPass(); + // render first clip path + beginRenderPass(commandEncoder, dst, true); + drawClipPath(context, (WgRenderDataShape*)(renderData->clips[0])); + endRenderPass(); + // render amd merge clip paths + for (uint32_t i = 1 ; i < renderData->clips.count; i++) { + // render clip path + beginRenderPass(commandEncoder, &storageInterm, true); + drawClipPath(context, (WgRenderDataShape*)(renderData->clips[i])); + endRenderPass(); + // merge masks + mergeMasks(commandEncoder, &storageInterm, dst); + } + // restore current render pass + beginRenderPass(commandEncoder, target, false); +} + + +void WgCompositor::renderShape(WgContext& context, WgRenderDataShape* renderData, BlendMethod blendMethod) +{ + assert(renderData); + assert(renderPassEncoder); + WgPipelineBlendType blentType = blendMethodToBlendType(blendMethod); + // apply clip path if neccessary + if (renderData->clips.count != 0) { + renderClipPath(context, renderData, &storageClipPath); + if (renderData->strokeFirst) { + composeStrokes(context, renderData, &storageClipPath, CompositeMethod::ClipPath); + composeShape(context, renderData, &storageClipPath, CompositeMethod::ClipPath); + } else { + composeShape(context, renderData, &storageClipPath, CompositeMethod::ClipPath); + composeStrokes(context, renderData, &storageClipPath, CompositeMethod::ClipPath); + } + // use custom blending + } else if (blentType == WgPipelineBlendType::Custom) { + if (renderData->strokeFirst) { + blendStrokes(context, renderData, blendMethod); + blendShape(context, renderData, blendMethod); + } else { + blendShape(context, renderData, blendMethod); + blendStrokes(context, renderData, blendMethod); + } + // use direct hardware blending + } else { + if (renderData->strokeFirst) { + drawStrokes(context, renderData, blentType); + drawShape(context, renderData, blentType); + } else { + drawShape(context, renderData, blentType); + drawStrokes(context, renderData, blentType); + } + } +} + + +void WgCompositor::renderImage(WgContext& context, WgRenderDataPicture* renderData, BlendMethod blendMethod) +{ + assert(renderData); + assert(renderPassEncoder); + WgPipelineBlendType blentType = blendMethodToBlendType(blendMethod); + // apply clip path if neccessary + if (renderData->clips.count != 0) { + renderClipPath(context, renderData, &storageClipPath); + composeImage(context, renderData, &storageClipPath, CompositeMethod::ClipPath); + // use custom blending + } else if (blentType == WgPipelineBlendType::Custom) + blendImage(context, renderData, blendMethod); + else // use direct hardware blending + drawImage(context, renderData, blentType); +}; + + +void WgCompositor::blendShape(WgContext& context, WgRenderDataShape* renderData, BlendMethod blendMethod) +{ + assert(renderData); + assert(commandEncoder); + assert(currentTarget); + // skip shape blending if shapes do not exist + if (renderData->renderSettingsShape.skip) return; + if (renderData->meshGroupShapes.meshes.count == 0) return; + // store current render pass + WgRenderStorage *target = currentTarget; + endRenderPass(); + // render into intermediate buffer + beginRenderPass(commandEncoder, &storageInterm, true); + drawShape(context, renderData, WgPipelineBlendType::Custom); + endRenderPass(); + // run blend + blend(commandEncoder, &storageInterm, target, 255, blendMethod, renderData->renderSettingsShape.rasterType); + // restore current render pass + beginRenderPass(commandEncoder, target, false); +} + +void WgCompositor::blendStrokes(WgContext& context, WgRenderDataShape* renderData, BlendMethod blendMethod) +{ + assert(renderData); + assert(commandEncoder); + assert(currentTarget); + // skip strokes blending if strokes do not exist + if (renderData->renderSettingsStroke.skip) return; + if (renderData->meshGroupStrokes.meshes.count == 0) return; + // store current render pass + WgRenderStorage *target = currentTarget; + endRenderPass(); + // render into intermediate buffer + beginRenderPass(commandEncoder, &storageInterm, true); + drawStrokes(context, renderData, WgPipelineBlendType::Custom); + endRenderPass(); + // run blend + blend(commandEncoder, &storageInterm, target, 255, blendMethod, renderData->renderSettingsStroke.rasterType); + // restore current render pass + beginRenderPass(commandEncoder, target, false); +}; + + +void WgCompositor::blendImage(WgContext& context, WgRenderDataPicture* renderData, BlendMethod blendMethod) +{ + assert(renderData); + assert(commandEncoder); + assert(currentTarget); + // store current render pass + WgRenderStorage *target = currentTarget; + endRenderPass(); + // render into intermediate buffer + beginRenderPass(commandEncoder, &storageInterm, true); + drawImage(context, renderData, WgPipelineBlendType::Custom); + endRenderPass(); + // run blend + blend(commandEncoder, &storageInterm, target, 255, blendMethod, WgRenderRasterType::Image); + // restore current render pass + beginRenderPass(commandEncoder, target, false); +}; + + +void WgCompositor::composeShape(WgContext& context, WgRenderDataShape* renderData, WgRenderStorage* mask, CompositeMethod composeMethod) +{ + assert(mask); + assert(renderData); + assert(commandEncoder); + assert(currentTarget); + // skip shape composing if shape do not exist + if (renderData->renderSettingsShape.skip) return; + if (renderData->meshGroupShapes.meshes.count == 0) return; + // store current render pass + WgRenderStorage *target = currentTarget; + endRenderPass(); + // render into intermediate buffer + beginRenderPass(commandEncoder, &storageInterm, true); + drawShape(context, renderData, WgPipelineBlendType::Custom); + endRenderPass(); + // run compisition + compose(commandEncoder, &storageInterm, mask, target, composeMethod); + // restore current render pass + beginRenderPass(commandEncoder, target, false); +} + + +void WgCompositor::composeStrokes(WgContext& context, WgRenderDataShape* renderData, WgRenderStorage* msk, CompositeMethod composeMethod) +{ + assert(msk); + assert(renderData); + assert(commandEncoder); + assert(currentTarget); + // skip shape composing if strokes do not exist + if (renderData->renderSettingsStroke.skip) return; + if (renderData->meshGroupStrokes.meshes.count == 0) return; + // store current render pass + WgRenderStorage *target = currentTarget; + endRenderPass(); + // render into intermediate buffer + beginRenderPass(commandEncoder, &storageInterm, true); + drawStrokes(context, renderData, WgPipelineBlendType::Custom); + endRenderPass(); + // run compisition + compose(commandEncoder, &storageInterm, msk, target, composeMethod); + // restore current render pass + beginRenderPass(commandEncoder, target, false); +} + + +void WgCompositor::composeImage(WgContext& context, WgRenderDataPicture* renderData, WgRenderStorage* msk, CompositeMethod composeMethod) +{ + assert(msk); + assert(renderData); + assert(commandEncoder); + assert(currentTarget); + // store current render pass + WgRenderStorage *target = currentTarget; + endRenderPass(); + // render into intermediate buffer + beginRenderPass(commandEncoder, &storageInterm, true); + drawImage(context, renderData, WgPipelineBlendType::Custom); + endRenderPass(); + // run compisition + compose(commandEncoder, &storageInterm, msk, target, composeMethod); + // restore current render pass + beginRenderPass(commandEncoder, target, false); +} + + +void WgCompositor::drawClipPath(WgContext& context, WgRenderDataShape* renderData) +{ + assert(renderData); + assert(renderPassEncoder); + assert(renderData->meshGroupShapes.meshes.count == renderData->meshGroupShapesBBox.meshes.count); + if ((renderData->viewport.w <= 0) || (renderData->viewport.h <= 0)) return; + wgpuRenderPassEncoderSetScissorRect(renderPassEncoder, renderData->viewport.x, renderData->viewport.y, renderData->viewport.w, renderData->viewport.h); + // setup stencil rules + WGPURenderPipeline stencilPipeline = (renderData->fillRule == FillRule::Winding) ? pipelines->winding : pipelines->evenodd; + wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 0); + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 0, bindGroupViewMat, 0, nullptr); + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 1, renderData->bindGroupPaint, 0, nullptr); + wgpuRenderPassEncoderSetPipeline(renderPassEncoder, stencilPipeline); + // draw to stencil (first pass) + for (uint32_t i = 0; i < renderData->meshGroupShapes.meshes.count; i++) + renderData->meshGroupShapes.meshes[i]->drawFan(context, renderPassEncoder); + // setup fill rules + wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 0); + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 0, bindGroupViewMat, 0, nullptr); + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 1, renderData->bindGroupPaint, 0, nullptr); + wgpuRenderPassEncoderSetPipeline(renderPassEncoder, pipelines->clipPath); + // draw to color (second pass) + renderData->meshDataBBox.drawFan(context, renderPassEncoder); +} + + +void WgCompositor::drawShape(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType) +{ + assert(renderData); + assert(renderPassEncoder); + assert(renderData->meshGroupShapes.meshes.count == renderData->meshGroupShapesBBox.meshes.count); + uint32_t blendTypeInd = (uint32_t)blendType; + if (renderData->renderSettingsShape.skip) return; + if ((renderData->viewport.w <= 0) || (renderData->viewport.h <= 0)) return; + wgpuRenderPassEncoderSetScissorRect(renderPassEncoder, renderData->viewport.x, renderData->viewport.y, renderData->viewport.w, renderData->viewport.h); + // setup stencil rules + WGPURenderPipeline stencilPipeline = (renderData->fillRule == FillRule::Winding) ? pipelines->winding : pipelines->evenodd; + wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 0); + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 0, bindGroupViewMat, 0, nullptr); + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 1, renderData->bindGroupPaint, 0, nullptr); + wgpuRenderPassEncoderSetPipeline(renderPassEncoder, stencilPipeline); + // draw to stencil (first pass) + for (uint32_t i = 0; i < renderData->meshGroupShapes.meshes.count; i++) + renderData->meshGroupShapes.meshes[i]->drawFan(context, renderPassEncoder); + // setup fill rules + wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 0); + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 0, bindGroupViewMat, 0, nullptr); + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 1, renderData->bindGroupPaint, 0, nullptr); + WgRenderSettings& settings = renderData->renderSettingsShape; + if (settings.fillType == WgRenderSettingsType::Solid) { + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 2, settings.bindGroupSolid, 0, nullptr); + wgpuRenderPassEncoderSetPipeline(renderPassEncoder, pipelines->solid[blendTypeInd]); + } else if (settings.fillType == WgRenderSettingsType::Linear) { + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 2, settings.bindGroupLinear, 0, nullptr); + wgpuRenderPassEncoderSetPipeline(renderPassEncoder, pipelines->linear[blendTypeInd]); + } else if (settings.fillType == WgRenderSettingsType::Radial) { + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 2, settings.bindGroupRadial, 0, nullptr); + wgpuRenderPassEncoderSetPipeline(renderPassEncoder, pipelines->radial[blendTypeInd]); + } + // draw to color (second pass) + renderData->meshDataBBox.drawFan(context, renderPassEncoder); +} + + +void WgCompositor::drawStrokes(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType) +{ + assert(renderData); + assert(renderPassEncoder); + assert(renderData->meshGroupStrokes.meshes.count == renderData->meshGroupStrokesBBox.meshes.count); + uint32_t blendTypeInd = (uint32_t)blendType; + if (renderData->renderSettingsStroke.skip) return; + if ((renderData->viewport.w <= 0) || (renderData->viewport.h <= 0)) return; + wgpuRenderPassEncoderSetScissorRect(renderPassEncoder, renderData->viewport.x, renderData->viewport.y, renderData->viewport.w, renderData->viewport.h); + // draw strokes to stencil (first pass) + for (uint32_t i = 0; i < renderData->meshGroupStrokes.meshes.count; i++) { + // setup stencil rules + wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 255); + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 0, bindGroupViewMat, 0, nullptr); + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 1, renderData->bindGroupPaint, 0, nullptr); + wgpuRenderPassEncoderSetPipeline(renderPassEncoder, pipelines->direct); + // draw to stencil (first pass) + renderData->meshGroupStrokes.meshes[i]->draw(context, renderPassEncoder); + // setup fill rules + wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 0); + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 0, bindGroupViewMat, 0, nullptr); + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 1, renderData->bindGroupPaint, 0, nullptr); + WgRenderSettings& settings = renderData->renderSettingsStroke; + if (settings.fillType == WgRenderSettingsType::Solid) { + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 2, settings.bindGroupSolid, 0, nullptr); + wgpuRenderPassEncoderSetPipeline(renderPassEncoder, pipelines->solid[blendTypeInd]); + } else if (settings.fillType == WgRenderSettingsType::Linear) { + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 2, settings.bindGroupLinear, 0, nullptr); + wgpuRenderPassEncoderSetPipeline(renderPassEncoder, pipelines->linear[blendTypeInd]); + } else if (settings.fillType == WgRenderSettingsType::Radial) { + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 2, settings.bindGroupRadial, 0, nullptr); + wgpuRenderPassEncoderSetPipeline(renderPassEncoder, pipelines->radial[blendTypeInd]); + } + // draw to color (second pass) + renderData->meshGroupStrokesBBox.meshes[i]->drawFan(context, renderPassEncoder); + } +} + + +void WgCompositor::drawImage(WgContext& context, WgRenderDataPicture* renderData, WgPipelineBlendType blendType) +{ + assert(renderData); + assert(renderPassEncoder); + uint32_t blendTypeInd = (uint32_t)blendType; + if ((renderData->viewport.w <= 0) || (renderData->viewport.h <= 0)) return; + wgpuRenderPassEncoderSetScissorRect(renderPassEncoder, renderData->viewport.x, renderData->viewport.y, renderData->viewport.w, renderData->viewport.h); + wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 0); + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 0, bindGroupViewMat, 0, nullptr); + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 1, renderData->bindGroupPaint, 0, nullptr); + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 2, renderData->bindGroupPicture, 0, nullptr); + wgpuRenderPassEncoderSetPipeline(renderPassEncoder, pipelines->image[blendTypeInd]); + renderData->meshData.drawImage(context, renderPassEncoder); +} + + +void WgCompositor::mergeMasks(WGPUCommandEncoder encoder, WgRenderStorage* mask0, WgRenderStorage* mask1) +{ + assert(mask0); + assert(mask1); + assert(!renderPassEncoder); + // copy dst storage to temporary read only storage + const WGPUImageCopyTexture texSrc { .texture = mask1->texture }; + const WGPUImageCopyTexture texDst { .texture = storageDstCopy.texture }; + const WGPUExtent3D copySize { .width = width, .height = height, .depthOrArrayLayers = 1 }; + wgpuCommandEncoderCopyTextureToTexture(encoder, &texSrc, &texDst, ©Size); + // execute compose shader + const WGPUComputePassDescriptor computePassDescriptor{}; + WGPUComputePassEncoder computePassEncoder = wgpuCommandEncoderBeginComputePass(encoder, &computePassDescriptor); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 0, mask0->bindGroupRead, 0, nullptr); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 1, storageDstCopy.bindGroupRead, 0, nullptr); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 2, mask1->bindGroupWrite, 0, nullptr); + wgpuComputePassEncoderSetPipeline(computePassEncoder, pipelines->mergeMasks); + wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, (width + 7) / 8, (height + 7) / 8, 1); + wgpuComputePassEncoderEnd(computePassEncoder); +} + + +void WgCompositor::blend(WGPUCommandEncoder encoder, WgRenderStorage* src, WgRenderStorage* dst, uint8_t opacity, BlendMethod blendMethod, WgRenderRasterType rasterType) +{ + assert(src); + assert(dst); + assert(!renderPassEncoder); + WGPUComputePipeline pipeline = pipelines->blendImage[(size_t)blendMethod]; + if (rasterType == WgRenderRasterType::Solid) pipeline = pipelines->blendSolid[(size_t)blendMethod]; + if (rasterType == WgRenderRasterType::Gradient) pipeline = pipelines->blendGradient[(size_t)blendMethod]; + // copy texture to texture + const WGPUImageCopyTexture texSrc { .texture = dst->texture }; + const WGPUImageCopyTexture texDst { .texture = storageDstCopy.texture }; + const WGPUExtent3D copySize { .width = width, .height = height, .depthOrArrayLayers = 1 }; + wgpuCommandEncoderCopyTextureToTexture(encoder, &texSrc, &texDst, ©Size); + // execute compose shader + const WGPUComputePassDescriptor computePassDescriptor{}; + WGPUComputePassEncoder computePassEncoder = wgpuCommandEncoderBeginComputePass(encoder, &computePassDescriptor); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 0, src->bindGroupRead, 0, nullptr); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 1, storageDstCopy.bindGroupRead, 0, nullptr); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 2, dst->bindGroupWrite, 0, nullptr); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 3, bindGroupOpacities[opacity], 0, nullptr); + wgpuComputePassEncoderSetPipeline(computePassEncoder, pipeline); + wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, (width + 7) / 8, (height + 7) / 8, 1); + wgpuComputePassEncoderEnd(computePassEncoder); +} + + +void WgCompositor::compose(WGPUCommandEncoder encoder, WgRenderStorage* src, WgRenderStorage* mask, WgRenderStorage* dst, CompositeMethod composeMethod) +{ + assert(src); + assert(mask); + assert(dst); + assert(!renderPassEncoder); + // copy dst storage to temporary read only storage + const WGPUImageCopyTexture texSrc { .texture = dst->texture }; + const WGPUImageCopyTexture texDst { .texture = storageDstCopy.texture }; + const WGPUExtent3D copySize { .width = width, .height = height, .depthOrArrayLayers = 1 }; + wgpuCommandEncoderCopyTextureToTexture(encoder, &texSrc, &texDst, ©Size); + // execute compose shader + const WGPUComputePassDescriptor computePassDescriptor{}; + WGPUComputePassEncoder computePassEncoder = wgpuCommandEncoderBeginComputePass(encoder, &computePassDescriptor); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 0, src->bindGroupRead, 0, nullptr); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 1, mask->bindGroupRead, 0, nullptr); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 2, storageDstCopy.bindGroupRead, 0, nullptr); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 3, dst->bindGroupWrite, 0, nullptr); + wgpuComputePassEncoderSetPipeline(computePassEncoder, pipelines->compose[(size_t)composeMethod]); + wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, (width + 7) / 8, (height + 7) / 8, 1); + wgpuComputePassEncoderEnd(computePassEncoder); +} diff --git a/src/renderer/wg_engine/tvgWgCompositor.h b/src/renderer/wg_engine/tvgWgCompositor.h new file mode 100755 index 00000000..b4d6a409 --- /dev/null +++ b/src/renderer/wg_engine/tvgWgCompositor.h @@ -0,0 +1,90 @@ +/* + * Copyright (c) 2024 the ThorVG project. All rights reserved. + + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ + +#ifndef _TVG_WG_COMPOSITOR_H_ +#define _TVG_WG_COMPOSITOR_H_ + +#include "tvgWgRenderTarget.h" +#include "tvgWgRenderData.h" + +struct WgCompose: public Compositor { + BlendMethod blend; +}; + +class WgCompositor { +private: + // pipelines (external handle, do not release) + WgPipelines* pipelines{}; + // global stencil buffer handles + WGPUTexture texStencil{}; + WGPUTextureView texViewStencil{}; + // global view matrix handles + WGPUBuffer bufferViewMat{}; + WGPUBindGroup bindGroupViewMat{}; + // opacity value pool + WGPUBuffer bufferOpacities[256]{}; + WGPUBindGroup bindGroupOpacities[256]{}; + // current render pass handles + WGPURenderPassEncoder renderPassEncoder{}; + WGPUCommandEncoder commandEncoder{}; + WgRenderStorage* currentTarget{}; + // intermediate render storages + WgRenderStorage storageInterm; + WgRenderStorage storageClipPath; + WgRenderStorage storageDstCopy; + + static WgPipelineBlendType blendMethodToBlendType(BlendMethod blendMethod); +public: + // render target dimensions + uint32_t width{}; + uint32_t height{}; +public: + void initialize(WgContext& context, uint32_t width, uint32_t height); + void release(WgContext& context); + + void beginRenderPass(WGPUCommandEncoder encoder, WgRenderStorage* target, bool clear); + void endRenderPass(); + + void renderClipPath(WgContext& context, WgRenderDataPaint* renderData, WgRenderStorage* dst); + + void renderShape(WgContext& context, WgRenderDataShape* renderData, BlendMethod blendMethod); + void renderImage(WgContext& context, WgRenderDataPicture* renderData, BlendMethod blendMethod); + + void blendShape(WgContext& context, WgRenderDataShape* renderData, BlendMethod blendMethod); + void blendStrokes(WgContext& context, WgRenderDataShape* renderData, BlendMethod blendMethod); + void blendImage(WgContext& context, WgRenderDataPicture* renderData, BlendMethod blendMethod); + + void composeShape(WgContext& context, WgRenderDataShape* renderData, WgRenderStorage* mask, CompositeMethod composeMethod); + void composeStrokes(WgContext& context, WgRenderDataShape* renderData, WgRenderStorage* mask, CompositeMethod composeMethod); + void composeImage(WgContext& context, WgRenderDataPicture* renderData, WgRenderStorage* mask, CompositeMethod composeMethod); + + void drawClipPath(WgContext& context, WgRenderDataShape* renderData); + void drawShape(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType); + void drawStrokes(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType); + void drawImage(WgContext& context, WgRenderDataPicture* renderData, WgPipelineBlendType blendType); + + void mergeMasks(WGPUCommandEncoder encoder, WgRenderStorage* mask0, WgRenderStorage* mask1); + void blend(WGPUCommandEncoder encoder, WgRenderStorage* src, WgRenderStorage* dst, uint8_t opacity, BlendMethod blendMethod, WgRenderRasterType rasterType); + void compose(WGPUCommandEncoder encoder, WgRenderStorage* src, WgRenderStorage* mask, WgRenderStorage* dst, CompositeMethod composeMethod); +}; + +#endif // _TVG_WG_COMPOSITOR_H_ diff --git a/src/renderer/wg_engine/tvgWgGeometry.cpp b/src/renderer/wg_engine/tvgWgGeometry.cpp old mode 100644 new mode 100755 diff --git a/src/renderer/wg_engine/tvgWgGeometry.h b/src/renderer/wg_engine/tvgWgGeometry.h old mode 100644 new mode 100755 diff --git a/src/renderer/wg_engine/tvgWgPipelines.cpp b/src/renderer/wg_engine/tvgWgPipelines.cpp old mode 100644 new mode 100755 index 1d87fec4..7709c0df --- a/src/renderer/wg_engine/tvgWgPipelines.cpp +++ b/src/renderer/wg_engine/tvgWgPipelines.cpp @@ -23,492 +23,359 @@ #include "tvgWgPipelines.h" #include "tvgWgShaderSrc.h" -#define ARRAY_ELEMENTS_COUNT(arr) sizeof(arr)/sizeof(arr[0]) - -//************************************************************************ -// graphics pipelines -//************************************************************************ - -void WgPipelineFillShapeWinding::initialize(WGPUDevice device) +WGPUShaderModule WgPipelines::createShaderModule(WGPUDevice device, const char* label, const char* code) { - // vertex attributes settings - WGPUVertexAttribute vertexAttributesPos = { WGPUVertexFormat_Float32x2, sizeof(float) * 0, 0 }; - WGPUVertexBufferLayout vertexBufferLayouts[] = { - makeVertexBufferLayout(&vertexAttributesPos, 1, sizeof(float) * 2) - }; - - // bind groups - WGPUBindGroupLayout bindGroupLayouts[] = { - WgBindGroupCanvas::getLayout(device), - WgBindGroupPaint::getLayout(device) - }; - - // stencil function - WGPUCompareFunction stencilFunctionFront = WGPUCompareFunction_Always; - WGPUStencilOperation stencilOperationFront = WGPUStencilOperation_IncrementWrap; - WGPUCompareFunction stencilFunctionBack = WGPUCompareFunction_Always; - WGPUStencilOperation stencilOperationBack = WGPUStencilOperation_DecrementWrap; - - // shader source and labels - auto shaderSource = cShaderSource_PipelineFill; - auto shaderLabel = "The shader fill"; - auto pipelineLabel = "The render pipeline fill shape winding"; - - // allocate all pipeline handles - allocate(device, WgPipelineBlendType::SrcOver, WGPUColorWriteMask_None, - vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts), - bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts), - stencilFunctionFront, stencilOperationFront, stencilFunctionBack, stencilOperationBack, - shaderSource, shaderLabel, pipelineLabel); + WGPUShaderModuleWGSLDescriptor shaderModuleWGSLDesc{}; + shaderModuleWGSLDesc.chain.sType = WGPUSType_ShaderModuleWGSLDescriptor; + shaderModuleWGSLDesc.code = code; + const WGPUShaderModuleDescriptor shaderModuleDesc { .nextInChain = &shaderModuleWGSLDesc.chain, .label = label }; + return wgpuDeviceCreateShaderModule(device, &shaderModuleDesc); } -void WgPipelineFillShapeEvenOdd::initialize(WGPUDevice device) +WGPUPipelineLayout WgPipelines::createPipelineLayout(WGPUDevice device, const WGPUBindGroupLayout* bindGroupLayouts, const uint32_t bindGroupLayoutsCount) { - // vertex attributes settings - WGPUVertexAttribute vertexAttributesPos = { WGPUVertexFormat_Float32x2, sizeof(float) * 0, 0 }; - WGPUVertexBufferLayout vertexBufferLayouts[] = { - makeVertexBufferLayout(&vertexAttributesPos, 1, sizeof(float) * 2) - }; - - // bind groups - WGPUBindGroupLayout bindGroupLayouts[] = { - WgBindGroupCanvas::getLayout(device), - WgBindGroupPaint::getLayout(device) - }; - - // stencil function - WGPUCompareFunction stencilFunctionFront = WGPUCompareFunction_Always; - WGPUStencilOperation stencilOperationFront = WGPUStencilOperation_Invert; - WGPUCompareFunction stencilFunctionBack = WGPUCompareFunction_Always; - WGPUStencilOperation stencilOperationBack = WGPUStencilOperation_Invert; - - // shader source and labels - auto shaderSource = cShaderSource_PipelineFill; - auto shaderLabel = "The shader fill"; - auto pipelineLabel = "The render pipeline fill shape Even Odd"; - - // allocate all pipeline handles - allocate(device, WgPipelineBlendType::SrcOver, WGPUColorWriteMask_None, - vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts), - bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts), - stencilFunctionFront, stencilOperationFront, stencilFunctionBack, stencilOperationBack, - shaderSource, shaderLabel, pipelineLabel); + const WGPUPipelineLayoutDescriptor pipelineLayoutDesc { .bindGroupLayoutCount = bindGroupLayoutsCount, .bindGroupLayouts = bindGroupLayouts }; + return wgpuDeviceCreatePipelineLayout(device, &pipelineLayoutDesc); } -void WgPipelineFillStroke::initialize(WGPUDevice device) +WGPURenderPipeline WgPipelines::createRenderPipeline( + WGPUDevice device, const char* pipelineLabel, + const WGPUShaderModule shaderModule, const WGPUPipelineLayout pipelineLayout, + const WGPUVertexBufferLayout *vertexBufferLayouts, const uint32_t vertexBufferLayoutsCount, + const WGPUColorWriteMaskFlags writeMask, + const WGPUCompareFunction stencilFunctionFrnt, const WGPUStencilOperation stencilOperationFrnt, + const WGPUCompareFunction stencilFunctionBack, const WGPUStencilOperation stencilOperationBack, + const WGPUPrimitiveState primitiveState, const WGPUMultisampleState multisampleState, const WGPUBlendState blendState) { - // vertex and buffers settings - WGPUVertexAttribute vertexAttributesPos = { WGPUVertexFormat_Float32x2, sizeof(float) * 0, 0 }; - WGPUVertexBufferLayout vertexBufferLayouts[] = { - makeVertexBufferLayout(&vertexAttributesPos, 1, sizeof(float) * 2) + const WGPUColorTargetState colorTargetState { .format = WGPUTextureFormat_RGBA8Unorm, .blend = &blendState, .writeMask = writeMask }; + const WGPUColorTargetState colorTargetStates[] { colorTargetState }; + const WGPUDepthStencilState depthStencilState { + .format = WGPUTextureFormat_Stencil8, .depthCompare = WGPUCompareFunction_Always, + .stencilFront = { .compare = stencilFunctionFrnt, .failOp = stencilOperationFrnt, .depthFailOp = stencilOperationFrnt, .passOp = stencilOperationFrnt }, + .stencilBack = { .compare = stencilFunctionBack, .failOp = stencilOperationBack, .depthFailOp = stencilOperationBack, .passOp = stencilOperationBack }, + .stencilReadMask = 0xFFFFFFFF, .stencilWriteMask = 0xFFFFFFFF }; - - // bind groups and layouts - WGPUBindGroupLayout bindGroupLayouts[] = { - WgBindGroupCanvas::getLayout(device), - WgBindGroupPaint::getLayout(device) + const WGPUVertexState vertexState { .module = shaderModule, .entryPoint = "vs_main", .bufferCount = vertexBufferLayoutsCount, .buffers = vertexBufferLayouts }; + const WGPUFragmentState fragmentState { .module = shaderModule, .entryPoint = "fs_main", .targetCount = 1, .targets = colorTargetStates }; + const WGPURenderPipelineDescriptor renderPipelineDesc { + .label = pipelineLabel, + .layout = pipelineLayout, + .vertex = vertexState, + .primitive = primitiveState, + .depthStencil = &depthStencilState, + .multisample = multisampleState, + .fragment = &fragmentState }; - - // stencil function - WGPUCompareFunction stencilFunction = WGPUCompareFunction_Always; - WGPUStencilOperation stencilOperation = WGPUStencilOperation_Replace; - - // shader source and labels - auto shaderSource = cShaderSource_PipelineFill; - auto shaderLabel = "The shader fill"; - auto pipelineLabel = "The render pipeline fill stroke"; - - // allocate all pipeline handles - allocate(device, WgPipelineBlendType::SrcOver, WGPUColorWriteMask_None, - vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts), - bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts), - stencilFunction, stencilOperation, stencilFunction, stencilOperation, - shaderSource, shaderLabel, pipelineLabel); + return wgpuDeviceCreateRenderPipeline(device, &renderPipelineDesc); } -void WgPipelineClipMask::initialize(WGPUDevice device) +WGPUComputePipeline WgPipelines::createComputePipeline( + WGPUDevice device, const char* pipelineLabel, + const WGPUShaderModule shaderModule, const WGPUPipelineLayout pipelineLayout) { - // vertex and buffers settings - WGPUVertexAttribute vertexAttributesPos = { WGPUVertexFormat_Float32x2, sizeof(float) * 0, 0 }; - WGPUVertexBufferLayout vertexBufferLayouts[] = { - makeVertexBufferLayout(&vertexAttributesPos, 1, sizeof(float) * 2) + const WGPUComputePipelineDescriptor computePipelineDesc{ + .label = pipelineLabel, + .layout = pipelineLayout, + .compute = { + .module = shaderModule, + .entryPoint = "cs_main" + } }; - - // bind groups - WGPUBindGroupLayout bindGroupLayouts[] = { - WgBindGroupCanvas::getLayout(device), - WgBindGroupPaint::getLayout(device) - }; - - // stencil function - WGPUCompareFunction stencilFunction = WGPUCompareFunction_NotEqual; - WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero; - - // shader source and labels - auto shaderSource = cShaderSource_PipelineFill; - auto shaderLabel = "The shader fill"; - auto pipelineLabel = "The render pipeline clip mask"; - - // allocate all pipeline handles - allocate(device, WgPipelineBlendType::SrcOver, WGPUColorWriteMask_All, - vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts), - bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts), - stencilFunction, stencilOperation, stencilFunction, stencilOperation, - shaderSource, shaderLabel, pipelineLabel); + return wgpuDeviceCreateComputePipeline(device, &computePipelineDesc); } -void WgPipelineSolid::initialize(WGPUDevice device, WgPipelineBlendType blendType) +void WgPipelines::releaseComputePipeline(WGPUComputePipeline& computePipeline) { - // vertex and buffers settings - WGPUVertexAttribute vertexAttributesPos = { WGPUVertexFormat_Float32x2, sizeof(float) * 0, 0 }; - WGPUVertexBufferLayout vertexBufferLayouts[] = { - makeVertexBufferLayout(&vertexAttributesPos, 1, sizeof(float) * 2) - }; - - // bind groups and layouts - WGPUBindGroupLayout bindGroupLayouts[] = { - WgBindGroupCanvas::getLayout(device), - WgBindGroupPaint::getLayout(device), - WgBindGroupSolidColor::getLayout(device) - }; - - // stencil function - WGPUCompareFunction stencilFunction = WGPUCompareFunction_NotEqual; - WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero; - - // shader source and labels - auto shaderSource = cShaderSource_PipelineSolid; - auto shaderLabel = "The shader solid color"; - auto pipelineLabel = "The render pipeline solid color"; - - // allocate all pipeline handles - allocate(device, blendType, WGPUColorWriteMask_All, - vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts), - bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts), - stencilFunction, stencilOperation, stencilFunction, stencilOperation, - shaderSource, shaderLabel, pipelineLabel); + if (computePipeline) { + wgpuComputePipelineRelease(computePipeline); + computePipeline = nullptr; + } } -void WgPipelineLinear::initialize(WGPUDevice device, WgPipelineBlendType blendType) +void WgPipelines::releaseRenderPipeline(WGPURenderPipeline& renderPipeline) { - // vertex and buffers settings - WGPUVertexAttribute vertexAttributesPos = { WGPUVertexFormat_Float32x2, sizeof(float) * 0, 0 }; - WGPUVertexBufferLayout vertexBufferLayouts[] = { - makeVertexBufferLayout(&vertexAttributesPos, 1, sizeof(float) * 2) - }; - - // bind groups and layouts - WGPUBindGroupLayout bindGroupLayouts[] = { - WgBindGroupCanvas::getLayout(device), - WgBindGroupPaint::getLayout(device), - WgBindGroupLinearGradient::getLayout(device) - }; - - // stencil function - WGPUCompareFunction stencilFunction = WGPUCompareFunction_NotEqual; - WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero; - - // shader source and labels - auto shaderSource = cShaderSource_PipelineLinear; - auto shaderLabel = "The shader linear gradient"; - auto pipelineLabel = "The render pipeline linear gradient"; - - // allocate all pipeline handles - allocate(device, blendType, WGPUColorWriteMask_All, - vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts), - bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts), - stencilFunction, stencilOperation, stencilFunction, stencilOperation, - shaderSource, shaderLabel, pipelineLabel); + if (renderPipeline) { + wgpuRenderPipelineRelease(renderPipeline); + renderPipeline = nullptr; + } } -void WgPipelineRadial::initialize(WGPUDevice device, WgPipelineBlendType blendType) +void WgPipelines::releasePipelineLayout(WGPUPipelineLayout& pipelineLayout) { - // vertex and buffers settings - WGPUVertexAttribute vertexAttributesPos = { WGPUVertexFormat_Float32x2, sizeof(float) * 0, 0 }; - WGPUVertexBufferLayout vertexBufferLayouts[] = { - makeVertexBufferLayout(&vertexAttributesPos, 1, sizeof(float) * 2) - }; - - // bind groups and layouts - WGPUBindGroupLayout bindGroupLayouts[] = { - WgBindGroupCanvas::getLayout(device), - WgBindGroupPaint::getLayout(device), - WgBindGroupRadialGradient::getLayout(device) - }; - - // stencil function - WGPUCompareFunction stencilFunction = WGPUCompareFunction_NotEqual; - WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero; - - // shader source and labels - auto shaderSource = cShaderSource_PipelineRadial; - auto shaderLabel = "The shader radial gradient"; - auto pipelineLabel = "The render pipeline radial gradient"; - - // allocate all pipeline handles - allocate(device, blendType, WGPUColorWriteMask_All, - vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts), - bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts), - stencilFunction, stencilOperation, stencilFunction, stencilOperation, - shaderSource, shaderLabel, pipelineLabel); + if (pipelineLayout) { + wgpuPipelineLayoutRelease(pipelineLayout); + pipelineLayout = nullptr; + } } -void WgPipelineImage::initialize(WGPUDevice device, WgPipelineBlendType blendType) +void WgPipelines::releaseShaderModule(WGPUShaderModule& shaderModule) { - // 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[] = { - WgBindGroupCanvas::getLayout(device), - WgBindGroupPaint::getLayout(device), - WgBindGroupPicture::getLayout(device) - }; - - // stencil function - WGPUCompareFunction stencilFunction = WGPUCompareFunction_Always; - WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero; - - // shader source and labels - auto shaderSource = cShaderSource_PipelineImage; - auto shaderLabel = "The shader image"; - auto pipelineLabel = "The render pipeline image"; - - // allocate all pipeline handles - allocate(device, blendType, WGPUColorWriteMask_All, - vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts), - bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts), - stencilFunction, stencilOperation, stencilFunction, stencilOperation, - shaderSource, shaderLabel, pipelineLabel); + if (shaderModule) { + wgpuShaderModuleRelease(shaderModule); + shaderModule = nullptr; + } } -//************************************************************************ -// compute pipelines -//************************************************************************ - -void WgPipelineCopy::initialize(WGPUDevice device) -{ - // bind groups and layouts - WGPUBindGroupLayout bindGroupLayouts[] = { - WgBindGroupTextureStorageRgbaRO::getLayout(device), - WgBindGroupTextureStorageRgbaWO::getLayout(device) - }; - - // shader source and labels - auto shaderSource = cShaderSource_PipelineComputeCopy; - auto shaderLabel = "The compute shader copy"; - auto pipelineLabel = "The compute pipeline copy"; - - // allocate all pipeline handles - allocate(device, - bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts), - shaderSource, shaderLabel, pipelineLabel); -} - - -void WgPipelineBlend::initialize(WGPUDevice device, const char *shaderSource) -{ - // bind groups and layouts - WGPUBindGroupLayout bindGroupLayouts[] = { - WgBindGroupTexBlend::getLayout(device), - WgBindGroupBlendMethod::getLayout(device), - WgBindGroupOpacity::getLayout(device) - }; - - // shader source and labels - 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 WgPipelineBlendMask::initialize(WGPUDevice device, const char *shaderSource) -{ - // bind groups and layouts - WGPUBindGroupLayout bindGroupLayouts[] = { - WgBindGroupTexBlendMask::getLayout(device), - WgBindGroupBlendMethod::getLayout(device), - WgBindGroupOpacity::getLayout(device) - }; - - // shader source and labels - auto shaderLabel = "The compute shader blend mask"; - auto pipelineLabel = "The compute pipeline blend mask"; - - // allocate all pipeline handles - allocate(device, - bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts), - shaderSource, shaderLabel, pipelineLabel); -} - - -void WgPipelineMaskCompose::initialize(WGPUDevice device) -{ - // bind groups and layouts - WGPUBindGroupLayout bindGroupLayouts[] = { - WgBindGroupTexMaskCompose::getLayout(device), - }; - - // shader source and labels - auto shaderSource = cShaderSource_PipelineComputeMaskCompose; - auto shaderLabel = "The compute shader mask compose"; - auto pipelineLabel = "The compute pipeline mask compose"; - - // allocate all pipeline handles - allocate(device, - bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts), - shaderSource, shaderLabel, pipelineLabel); -} - - -void WgPipelineCompose::initialize(WGPUDevice device) -{ - // bind groups and layouts - WGPUBindGroupLayout bindGroupLayouts[] = { - WgBindGroupTexCompose::getLayout(device), - WgBindGroupCompositeMethod::getLayout(device), - WgBindGroupBlendMethod::getLayout(device), - WgBindGroupOpacity::getLayout(device) - }; - - // shader source and labels - auto shaderSource = cShaderSource_PipelineComputeCompose; - auto shaderLabel = "The compute shader compose blend"; - auto pipelineLabel = "The compute pipeline compose blend"; - - // allocate all pipeline handles - allocate(device, - bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts), - shaderSource, shaderLabel, pipelineLabel); -} - - -void WgPipelineAntiAliasing::initialize(WGPUDevice device) -{ - // bind groups and layouts - WGPUBindGroupLayout bindGroupLayouts[] = { - WgBindGroupTextureStorageRgbaRO::getLayout(device), - WgBindGroupTextureStorageBgraWO::getLayout(device) - }; - - // shader source and labels - auto shaderSource = cShaderSource_PipelineComputeAntiAlias; - auto shaderLabel = "The compute shader anti-aliasing"; - auto pipelineLabel = "The compute pipeline anti-aliasing"; - - // allocate all pipeline handles - allocate(device, - bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts), - shaderSource, shaderLabel, pipelineLabel); -} - -//************************************************************************ -// pipelines -//************************************************************************ void WgPipelines::initialize(WgContext& context) { - // fill pipelines - fillShapeWinding.initialize(context.device); - fillShapeEvenOdd.initialize(context.device); - fillStroke.initialize(context.device); - clipMask.initialize(context.device); - for (uint8_t type = (uint8_t)WgPipelineBlendType::SrcOver; type <= (uint8_t)WgPipelineBlendType::Custom; type++) { - solid[type].initialize(context.device, (WgPipelineBlendType)type); - linear[type].initialize(context.device, (WgPipelineBlendType)type); - radial[type].initialize(context.device, (WgPipelineBlendType)type); - image[type].initialize(context.device, (WgPipelineBlendType)type); - } - // compute pipelines - computeCopy.initialize(context.device); - computeBlendSolid.initialize(context.device, cShaderSource_PipelineComputeBlendSolid); - computeBlendGradient.initialize(context.device, cShaderSource_PipelineComputeBlendGradient); - computeBlendImage.initialize(context.device, cShaderSource_PipelineComputeBlendImage); - computeBlendSolidMask.initialize(context.device, cShaderSource_PipelineComputeBlendSolidMask); - computeBlendGradientMask.initialize(context.device, cShaderSource_PipelineComputeBlendGradientMask); - computeBlendImageMask.initialize(context.device, cShaderSource_PipelineComputeBlendImageMask); - computeMaskCompose.initialize(context.device); - computeCompose.initialize(context.device); - computeAntiAliasing.initialize(context.device); - // store pipelines to context + // share pipelines to context + assert(!context.pipelines); context.pipelines = this; -} + // initialize bind group layouts + layouts.initialize(context); -void WgPipelines::release() -{ - WgBindGroupTexCompose::releaseLayout(); - WgBindGroupTexMaskCompose::releaseLayout(); - WgBindGroupTexBlendMask::releaseLayout(); - WgBindGroupTexBlend::releaseLayout(); - WgBindGroupTextureStorageBgraRO::releaseLayout(); - WgBindGroupTextureStorageBgraWO::releaseLayout(); - WgBindGroupTextureStorageRgbaRO::releaseLayout(); - WgBindGroupTextureStorageRgbaWO::releaseLayout(); - WgBindGroupTexture::releaseLayout(); - WgBindGroupOpacity::releaseLayout(); - WgBindGroupPicture::releaseLayout(); - WgBindGroupRadialGradient::releaseLayout(); - WgBindGroupLinearGradient::releaseLayout(); - WgBindGroupSolidColor::releaseLayout(); - WgBindGroupPaint::releaseLayout(); - WgBindGroupCanvas::releaseLayout(); - // compute pipelines - computeAntiAliasing.release(); - computeCompose.release(); - computeBlendImageMask.release(); - computeBlendGradientMask.release(); - computeBlendSolidMask.release(); - computeBlendImage.release(); - computeBlendGradient.release(); - computeBlendSolid.release(); - computeCopy.release(); - // fill pipelines - for (uint8_t type = (uint8_t)WgPipelineBlendType::SrcOver; type <= (uint8_t)WgPipelineBlendType::Custom; type++) { - image[type].release(); - radial[type].release(); - linear[type].release(); - solid[type].release(); + // common pipeline settings + const WGPUVertexAttribute vertexAttributePos { .format = WGPUVertexFormat_Float32x2, .offset = 0, .shaderLocation = 0 }; + const WGPUVertexAttribute vertexAttributeTex { .format = WGPUVertexFormat_Float32x2, .offset = 0, .shaderLocation = 1 }; + const WGPUVertexAttribute vertexAttributesPos[] { vertexAttributePos }; + const WGPUVertexAttribute vertexAttributesTex[] { vertexAttributeTex }; + const WGPUVertexBufferLayout vertexBufferLayoutPos { .arrayStride = 8, .stepMode = WGPUVertexStepMode_Vertex, .attributeCount = 1, .attributes = vertexAttributesPos }; + const WGPUVertexBufferLayout vertexBufferLayoutTex { .arrayStride = 8, .stepMode = WGPUVertexStepMode_Vertex, .attributeCount = 1, .attributes = vertexAttributesTex }; + const WGPUVertexBufferLayout vertexBufferLayoutsShape[] { vertexBufferLayoutPos }; + const WGPUVertexBufferLayout vertexBufferLayoutsImage[] { vertexBufferLayoutPos, vertexBufferLayoutTex }; + const WGPUPrimitiveState primitiveState { .topology = WGPUPrimitiveTopology_TriangleList }; + const WGPUMultisampleState multisampleState { .count = 1, .mask = 0xFFFFFFFF, .alphaToCoverageEnabled = false }; + + // blend states + const WGPUBlendState blendStateSrc { + .color = { .operation = WGPUBlendOperation_Add, .srcFactor = WGPUBlendFactor_One, .dstFactor = WGPUBlendFactor_Zero }, + .alpha = { .operation = WGPUBlendOperation_Add, .srcFactor = WGPUBlendFactor_One, .dstFactor = WGPUBlendFactor_Zero } + }; + const WGPUBlendState blendStateNrm { + .color = { .operation = WGPUBlendOperation_Add, .srcFactor = WGPUBlendFactor_One, .dstFactor = WGPUBlendFactor_OneMinusSrcAlpha }, + .alpha = { .operation = WGPUBlendOperation_Add, .srcFactor = WGPUBlendFactor_One, .dstFactor = WGPUBlendFactor_OneMinusSrcAlpha } + }; + const WGPUBlendState blendStates[] { + blendStateSrc, // WgPipelineBlendType::SrcOver + blendStateNrm, // WgPipelineBlendType::Normal + blendStateSrc // WgPipelineBlendType::Custom (same as SrcOver) + }; + + // bind group layouts + const WGPUBindGroupLayout bindGroupLayoutsStencil[] = { + layouts.layoutBuffer1Un, + layouts.layoutBuffer2Un + }; + const WGPUBindGroupLayout bindGroupLayoutsFill[] = { + layouts.layoutBuffer1Un, + layouts.layoutBuffer2Un, + layouts.layoutBuffer1Un + }; + const WGPUBindGroupLayout bindGroupLayoutsImage[] = { + layouts.layoutBuffer1Un, + layouts.layoutBuffer2Un, + layouts.layoutTexSampled + }; + const WGPUBindGroupLayout bindGroupLayoutsMergeMasks[] = { + layouts.layoutTexStrorage1RO, + layouts.layoutTexStrorage1RO, + layouts.layoutTexStrorage1WO + }; + const WGPUBindGroupLayout bindGroupLayoutsBlend[] = { + layouts.layoutTexStrorage1RO, + layouts.layoutTexStrorage1RO, + layouts.layoutTexStrorage1WO, + layouts.layoutBuffer1Un + }; + const WGPUBindGroupLayout bindGroupLayoutsCompose[] = { + layouts.layoutTexStrorage1RO, + layouts.layoutTexStrorage1RO, + layouts.layoutTexStrorage1RO, + layouts.layoutTexStrorage1WO + }; + const WGPUBindGroupLayout bindGroupLayoutsCopy[] = { + layouts.layoutTexStrorage1RO, + layouts.layoutTexScreen1WO + }; + + // pipeline layouts + layoutStencil = createPipelineLayout(context.device, bindGroupLayoutsStencil, 2); + layoutFill = createPipelineLayout(context.device, bindGroupLayoutsFill, 3); + layoutImage = createPipelineLayout(context.device, bindGroupLayoutsImage, 3); + layoutBlend = createPipelineLayout(context.device, bindGroupLayoutsBlend, 4); + layoutCompose = createPipelineLayout(context.device, bindGroupLayoutsCompose, 4); + layoutMergeMasks = createPipelineLayout(context.device, bindGroupLayoutsMergeMasks, 3); + layoutCopy = createPipelineLayout(context.device, bindGroupLayoutsCopy, 2); + + // graphics shader modules + shaderStencil = createShaderModule(context.device, "The shader stencil", cShaderSrc_Stencil); + shaderSolid = createShaderModule(context.device, "The shader solid", cShaderSrc_Solid); + shaderRadial = createShaderModule(context.device, "The shader radial", cShaderSrc_Radial); + shaderLinear = createShaderModule(context.device, "The shader linear", cShaderSrc_Linear); + shaderImage = createShaderModule(context.device, "The shader image", cShaderSrc_Image); + // computes shader modules + shaderMergeMasks = createShaderModule(context.device, "The shader merge mask", cShaderSrc_MergeMasks); + shaderCopy = createShaderModule(context.device, "The shader copy", cShaderSrc_Copy); + + // render pipeline winding + winding = createRenderPipeline( + context.device, "The render pipeline winding", + shaderStencil, layoutStencil, + vertexBufferLayoutsShape, 1, + WGPUColorWriteMask_None, + WGPUCompareFunction_Always, WGPUStencilOperation_IncrementWrap, + WGPUCompareFunction_Always, WGPUStencilOperation_DecrementWrap, + primitiveState, multisampleState, blendStateSrc); + // render pipeline even-odd + evenodd = createRenderPipeline( + context.device, "The render pipeline even-odd", + shaderStencil, layoutStencil, + vertexBufferLayoutsShape, 1, + WGPUColorWriteMask_None, + WGPUCompareFunction_Always, WGPUStencilOperation_Invert, + WGPUCompareFunction_Always, WGPUStencilOperation_Invert, + primitiveState, multisampleState, blendStateSrc); + // render pipeline direct + direct = createRenderPipeline( + context.device, "The render pipeline direct", + shaderStencil, layoutStencil, + vertexBufferLayoutsShape, 1, + WGPUColorWriteMask_None, + WGPUCompareFunction_Always, WGPUStencilOperation_Replace, + WGPUCompareFunction_Always, WGPUStencilOperation_Replace, + primitiveState, multisampleState, blendStateSrc); + // render pipeline clip path + clipPath = createRenderPipeline( + context.device, "The render pipeline clip path", + shaderStencil, layoutStencil, + vertexBufferLayoutsShape, 1, + WGPUColorWriteMask_All, + WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero, + WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero, + primitiveState, multisampleState, blendStateSrc); + + // render pipeline solid + for (uint32_t i = 0; i < 3; i++) { + solid[i] = createRenderPipeline( + context.device, "The render pipeline solid", + shaderSolid, layoutFill, + vertexBufferLayoutsShape, 1, + WGPUColorWriteMask_All, + WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero, + WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero, + primitiveState, multisampleState, blendStates[i]); + } + + // render pipeline radial + for (uint32_t i = 0; i < 3; i++) { + radial[i] = createRenderPipeline( + context.device, "The render pipeline radial", + shaderRadial, layoutFill, + vertexBufferLayoutsShape, 1, + WGPUColorWriteMask_All, + WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero, + WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero, + primitiveState, multisampleState, blendStates[i]); + } + + // render pipeline linear + for (uint32_t i = 0; i < 3; i++) { + linear[i] = createRenderPipeline( + context.device, "The render pipeline linear", + shaderLinear, layoutFill, + vertexBufferLayoutsShape, 1, + WGPUColorWriteMask_All, + WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero, + WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero, + primitiveState, multisampleState, blendStates[i]); + } + + // render pipeline image + for (uint32_t i = 0; i < 3; i++) { + image[i] = createRenderPipeline( + context.device, "The render pipeline image", + shaderImage, layoutImage, + vertexBufferLayoutsImage, 2, + WGPUColorWriteMask_All, + WGPUCompareFunction_Always, WGPUStencilOperation_Zero, + WGPUCompareFunction_Always, WGPUStencilOperation_Zero, + primitiveState, multisampleState, blendStates[i]); + } + + // compute pipelines + mergeMasks = createComputePipeline(context.device, "The pipeline merge masks", shaderMergeMasks, layoutMergeMasks); + copy = createComputePipeline(context.device, "The pipeline copy", shaderCopy, layoutCopy); + // compute pipelines blend + const size_t shaderBlendCnt = sizeof(cShaderSrc_Blend_Solid)/sizeof(cShaderSrc_Blend_Solid[0]); + for (uint32_t i = 0; i < shaderBlendCnt; i++) { + // blend shaders + shaderBlendSolid[i] = createShaderModule(context.device, "The shader blend solid", cShaderSrc_Blend_Solid[i]); + shaderBlendGradient[i] = createShaderModule(context.device, "The shader blend gradient", cShaderSrc_Blend_Gradient[i]); + shaderBlendImage[i] = createShaderModule(context.device, "The shader blend image", cShaderSrc_Blend_Image[i]); + // blend pipelines + blendSolid[i] = createComputePipeline(context.device, "The pipeline blend solid", shaderBlendSolid[i], layoutBlend); + blendGradient[i] = createComputePipeline(context.device, "The pipeline blend gradient", shaderBlendGradient[i], layoutBlend); + blendImage[i] = createComputePipeline(context.device, "The pipeline blend image", shaderBlendImage[i], layoutBlend); + } + // compute pipelines compose + const size_t shaderComposeCnt = sizeof(cShaderSrc_Compose)/sizeof(cShaderSrc_Compose[0]); + for (uint32_t i = 0; i < shaderComposeCnt; i++) { + shaderCompose[i] = createShaderModule(context.device, "The shader compose", cShaderSrc_Compose[i]); + compose[i] = createComputePipeline(context.device, "The pipeline compose", shaderCompose[i], layoutCompose); } - clipMask.release(); - fillStroke.release(); - fillShapeEvenOdd.release(); - fillShapeWinding.release(); } - -bool WgPipelines::isBlendMethodSupportsHW(BlendMethod blendMethod) +void WgPipelines::releaseGraphicHandles(WgContext& context) { - switch (blendMethod) { - case BlendMethod::SrcOver: - case BlendMethod::Normal: - return true; - default: return false; - }; + releaseRenderPipeline(clipPath); + for (uint32_t i = 0; i < 3; i++) { + releaseRenderPipeline(image[i]); + releaseRenderPipeline(linear[i]); + releaseRenderPipeline(radial[i]); + releaseRenderPipeline(solid[i]); + } + releaseRenderPipeline(direct); + releaseRenderPipeline(evenodd); + releaseRenderPipeline(winding); + releasePipelineLayout(layoutImage); + releasePipelineLayout(layoutFill); + releasePipelineLayout(layoutStencil); + releaseShaderModule(shaderImage); + releaseShaderModule(shaderLinear); + releaseShaderModule(shaderRadial); + releaseShaderModule(shaderSolid); + releaseShaderModule(shaderStencil); } -WgPipelineBlendType WgPipelines::blendMethodToBlendType(BlendMethod blendMethod) +void WgPipelines::releaseComputeHandles(WgContext& context) { - switch (blendMethod) { - case BlendMethod::SrcOver: return WgPipelineBlendType::SrcOver; - case BlendMethod::Normal: return WgPipelineBlendType::Normal; - default: return WgPipelineBlendType::Custom; - }; + const size_t shaderComposeCnt = sizeof(shaderCompose)/sizeof(shaderCompose[0]); + for (uint32_t i = 0; i < shaderComposeCnt; i++) { + releaseComputePipeline(compose[i]); + releaseShaderModule(shaderCompose[i]); + } + const size_t shaderBlendImageCnt = sizeof(shaderBlendImage)/sizeof(shaderBlendImage[0]); + for (uint32_t i = 0; i < shaderBlendImageCnt; i++) { + releaseComputePipeline(blendImage[i]); + releaseComputePipeline(blendSolid[i]); + releaseComputePipeline(blendGradient[i]); + releaseShaderModule(shaderBlendImage[i]); + releaseShaderModule(shaderBlendGradient[i]); + releaseShaderModule(shaderBlendSolid[i]); + } + releaseComputePipeline(copy); + releaseComputePipeline(mergeMasks); + releasePipelineLayout(layoutCompose); + releasePipelineLayout(layoutBlend); + releasePipelineLayout(layoutMergeMasks); + releasePipelineLayout(layoutCopy); + releaseShaderModule(shaderMergeMasks); + releaseShaderModule(shaderCopy); +} + +void WgPipelines::release(WgContext& context) +{ + releaseComputeHandles(context); + releaseGraphicHandles(context); + layouts.release(context); } diff --git a/src/renderer/wg_engine/tvgWgPipelines.h b/src/renderer/wg_engine/tvgWgPipelines.h old mode 100644 new mode 100755 index 092cd6f5..69bcf8d4 --- a/src/renderer/wg_engine/tvgWgPipelines.h +++ b/src/renderer/wg_engine/tvgWgPipelines.h @@ -25,219 +25,75 @@ #include "tvgWgBindGroups.h" -//***************************************************************************** -// render pipelines -//***************************************************************************** - -struct WgPipelineFillShapeWinding: public WgRenderPipeline -{ - void initialize(WGPUDevice device) override; - void use(WGPURenderPassEncoder encoder, WgBindGroupCanvas& groupCanvas, WgBindGroupPaint& groupPaint) - { - set(encoder); - groupCanvas.set(encoder, 0); - groupPaint.set(encoder, 1); - } -}; - -struct WgPipelineFillShapeEvenOdd: public WgRenderPipeline -{ - void initialize(WGPUDevice device) override; - void use(WGPURenderPassEncoder encoder, WgBindGroupCanvas& groupCanvas, WgBindGroupPaint& groupPaint) - { - set(encoder); - groupCanvas.set(encoder, 0); - groupPaint.set(encoder, 1); - } -}; - -struct WgPipelineFillStroke: public WgRenderPipeline -{ - void initialize(WGPUDevice device) override; - void use(WGPURenderPassEncoder encoder, WgBindGroupCanvas& groupCanvas, WgBindGroupPaint& groupPaint) - { - set(encoder); - groupCanvas.set(encoder, 0); - groupPaint.set(encoder, 1); - } -}; - -struct WgPipelineClipMask: public WgRenderPipeline -{ - void initialize(WGPUDevice device) override; - void use(WGPURenderPassEncoder encoder, WgBindGroupCanvas& groupCanvas,WgBindGroupPaint& groupPaint) - { - set(encoder); - groupCanvas.set(encoder, 0); - groupPaint.set(encoder, 1); - } -}; - -struct WgPipelineSolid: public WgRenderPipeline -{ - void initialize(WGPUDevice device) override {} - void initialize(WGPUDevice device, WgPipelineBlendType blendType); - void use(WGPURenderPassEncoder encoder, WgBindGroupCanvas& groupCanvas,WgBindGroupPaint& groupPaint, WgBindGroupSolidColor& groupSolid) - { - set(encoder); - groupCanvas.set(encoder, 0); - groupPaint.set(encoder, 1); - groupSolid.set(encoder, 2); - } -}; - -struct WgPipelineLinear: public WgRenderPipeline -{ - void initialize(WGPUDevice device) override {} - void initialize(WGPUDevice device, WgPipelineBlendType blendType); - void use(WGPURenderPassEncoder encoder, WgBindGroupCanvas& groupCanvas, WgBindGroupPaint& groupPaint, WgBindGroupLinearGradient& groupLinear) - { - set(encoder); - groupCanvas.set(encoder, 0); - groupPaint.set(encoder, 1); - groupLinear.set(encoder, 2); - } -}; - -struct WgPipelineRadial: public WgRenderPipeline -{ - void initialize(WGPUDevice device) override {} - void initialize(WGPUDevice device, WgPipelineBlendType blendType); - void use(WGPURenderPassEncoder encoder, WgBindGroupCanvas& groupCanvas, WgBindGroupPaint& groupPaint, WgBindGroupRadialGradient& groupRadial) - { - set(encoder); - groupCanvas.set(encoder, 0); - groupPaint.set(encoder, 1); - groupRadial.set(encoder, 2); - } -}; - -struct WgPipelineImage: public WgRenderPipeline -{ - void initialize(WGPUDevice device) override { assert(false); }; - void initialize(WGPUDevice device, WgPipelineBlendType blendType); - void use(WGPURenderPassEncoder encoder, WgBindGroupCanvas& groupCanvas, WgBindGroupPaint& groupPaint, WgBindGroupPicture& groupPicture) - { - set(encoder); - groupCanvas.set(encoder, 0); - groupPaint.set(encoder, 1); - groupPicture.set(encoder, 2); - } -}; - -//***************************************************************************** -// compute pipelines -//***************************************************************************** - -struct WgPipelineCopy: public WgComputePipeline -{ - void initialize(WGPUDevice device) override; - void use(WGPUComputePassEncoder encoder, WgBindGroupTextureStorageRgbaRO& groupTexSrc, WgBindGroupTextureStorageRgbaWO& groupTexDst) - { - set(encoder); - groupTexSrc.set(encoder, 0); - groupTexDst.set(encoder, 1); - } -}; - - -struct WgPipelineBlend: public WgComputePipeline -{ - void initialize(WGPUDevice device) override { assert(false); }; - void initialize(WGPUDevice device, const char *shaderSource); - void use(WGPUComputePassEncoder encoder, WgBindGroupTexBlend& groupTexBlend, WgBindGroupBlendMethod& blendMethod, WgBindGroupOpacity& groupOpacity) - { - set(encoder); - groupTexBlend.set(encoder, 0); - blendMethod.set(encoder, 1); - groupOpacity.set(encoder, 2); - } -}; - - -struct WgPipelineBlendMask: public WgComputePipeline -{ - void initialize(WGPUDevice device) override { assert(false); }; - void initialize(WGPUDevice device, const char *shaderSource); - void use(WGPUComputePassEncoder encoder, WgBindGroupTexBlendMask& groupTexBlendMask, WgBindGroupBlendMethod& blendMethod, WgBindGroupOpacity& groupOpacity) - { - set(encoder); - groupTexBlendMask.set(encoder, 0); - blendMethod.set(encoder, 1); - groupOpacity.set(encoder, 2); - } -}; - - -struct WgPipelineMaskCompose: public WgComputePipeline -{ - void initialize(WGPUDevice device) override; - void use(WGPUComputePassEncoder encoder, WgBindGroupTexMaskCompose& groupTexMaskCompose) - { - set(encoder); - groupTexMaskCompose.set(encoder, 0); - } -}; - - -struct WgPipelineCompose: public WgComputePipeline -{ - void initialize(WGPUDevice device) override; - void use(WGPUComputePassEncoder encoder, WgBindGroupTexCompose& groupTexCompose, WgBindGroupCompositeMethod& groupComposeMethod, WgBindGroupBlendMethod& groupBlendMethod, WgBindGroupOpacity& groupOpacity) - { - set(encoder); - groupTexCompose.set(encoder, 0); - groupComposeMethod.set(encoder, 1); - groupBlendMethod.set(encoder, 2); - groupOpacity.set(encoder, 3); - } -}; - - -struct WgPipelineAntiAliasing: public WgComputePipeline -{ - void initialize(WGPUDevice device) override; - void use(WGPUComputePassEncoder encoder, WgBindGroupTextureStorageRgbaRO& groupTexSrc, WgBindGroupTextureStorageBgraWO& groupTexDst) - { - set(encoder); - groupTexSrc.set(encoder, 0); - groupTexDst.set(encoder, 1); - } -}; - -//***************************************************************************** -// pipelines -//***************************************************************************** - -struct WgPipelines -{ - // render pipelines - WgPipelineFillShapeWinding fillShapeWinding; - WgPipelineFillShapeEvenOdd fillShapeEvenOdd; - WgPipelineFillStroke fillStroke; - // fill pipelines - WgPipelineClipMask clipMask; - WgPipelineSolid solid[3]; - WgPipelineLinear linear[3]; - WgPipelineRadial radial[3]; - WgPipelineImage image[3]; - // compute pipelines - WgPipelineCopy computeCopy; - WgPipelineBlend computeBlendSolid; - WgPipelineBlend computeBlendGradient; - WgPipelineBlend computeBlendImage; - WgPipelineBlendMask computeBlendSolidMask; - WgPipelineBlendMask computeBlendGradientMask; - WgPipelineBlendMask computeBlendImageMask; - WgPipelineMaskCompose computeMaskCompose; - WgPipelineCompose computeCompose; - WgPipelineAntiAliasing computeAntiAliasing; +enum class WgPipelineBlendType { SrcOver = 0, Normal, Custom }; +class WgPipelines { +private: + // graphics pipeline shaders + WGPUShaderModule shaderStencil{}; + WGPUShaderModule shaderSolid{}; + WGPUShaderModule shaderRadial{}; + WGPUShaderModule shaderLinear{}; + WGPUShaderModule shaderImage{}; + // compute pipeline shaders + WGPUShaderModule shaderMergeMasks; + WGPUShaderModule shaderBlendSolid[14]; + WGPUShaderModule shaderBlendGradient[14]; + WGPUShaderModule shaderBlendImage[14]; + WGPUShaderModule shaderCompose[10]; + WGPUShaderModule shaderCopy; +private: + // graphics pipeline layouts + WGPUPipelineLayout layoutStencil{}; + WGPUPipelineLayout layoutFill{}; + WGPUPipelineLayout layoutImage{}; + // compute pipeline layouts + WGPUPipelineLayout layoutMergeMasks{}; + WGPUPipelineLayout layoutBlend{}; + WGPUPipelineLayout layoutCompose{}; + WGPUPipelineLayout layoutCopy{}; +public: + // graphics pipeline + WgBindGroupLayouts layouts; + WGPURenderPipeline winding{}; + WGPURenderPipeline evenodd{}; + WGPURenderPipeline direct{}; + WGPURenderPipeline solid[3]{}; + WGPURenderPipeline radial[3]{}; + WGPURenderPipeline linear[3]{}; + WGPURenderPipeline image[3]{}; + WGPURenderPipeline clipPath{}; + // compute pipeline + WGPUComputePipeline mergeMasks; + WGPUComputePipeline blendSolid[14]; + WGPUComputePipeline blendGradient[14]; + WGPUComputePipeline blendImage[14]; + WGPUComputePipeline compose[10]; + WGPUComputePipeline copy; +private: + void releaseGraphicHandles(WgContext& context); + void releaseComputeHandles(WgContext& context); +private: + WGPUShaderModule createShaderModule(WGPUDevice device, const char* label, const char* code); + WGPUPipelineLayout createPipelineLayout(WGPUDevice device, const WGPUBindGroupLayout* bindGroupLayouts, const uint32_t bindGroupLayoutsCount); + WGPURenderPipeline createRenderPipeline( + WGPUDevice device, const char* pipelineLabel, + const WGPUShaderModule shaderModule, const WGPUPipelineLayout pipelineLayout, + const WGPUVertexBufferLayout *vertexBufferLayouts, const uint32_t vertexBufferLayoutsCount, + const WGPUColorWriteMaskFlags writeMask, + const WGPUCompareFunction stencilFunctionFrnt, const WGPUStencilOperation stencilOperationFrnt, + const WGPUCompareFunction stencilFunctionBack, const WGPUStencilOperation stencilOperationBack, + const WGPUPrimitiveState primitiveState, const WGPUMultisampleState multisampleState, const WGPUBlendState blendState); + WGPUComputePipeline createComputePipeline( + WGPUDevice device, const char* pipelineLabel, + const WGPUShaderModule shaderModule, const WGPUPipelineLayout pipelineLayout); + void releaseComputePipeline(WGPUComputePipeline& computePipeline); + void releaseRenderPipeline(WGPURenderPipeline& renderPipeline); + void releasePipelineLayout(WGPUPipelineLayout& pipelineLayout); + void releaseShaderModule(WGPUShaderModule& shaderModule); +public: void initialize(WgContext& context); - void release(); - - static bool isBlendMethodSupportsHW(BlendMethod blendMethod); - static WgPipelineBlendType blendMethodToBlendType(BlendMethod blendMethod); + void release(WgContext& context); }; #endif // _TVG_WG_PIPELINES_H_ diff --git a/src/renderer/wg_engine/tvgWgRenderData.cpp b/src/renderer/wg_engine/tvgWgRenderData.cpp old mode 100644 new mode 100755 index 8dfc3a19..9e7c9676 --- a/src/renderer/wg_engine/tvgWgRenderData.cpp +++ b/src/renderer/wg_engine/tvgWgRenderData.cpp @@ -21,11 +21,9 @@ * SOFTWARE. */ -#ifndef _TVG_WG_RENDER_DATA_H_ -#define _TVG_WG_RENDER_DATA_H_ - #include #include "tvgWgRenderData.h" +#include "tvgWgShaderTypes.h" //*********************************************************************** // WgMeshData @@ -42,7 +40,7 @@ void WgMeshData::draw(WgContext& context, WGPURenderPassEncoder renderPassEncode void WgMeshData::drawFan(WgContext& context, WGPURenderPassEncoder renderPassEncoder) { wgpuRenderPassEncoderSetVertexBuffer(renderPassEncoder, 0, bufferPosition, 0, vertexCount * sizeof(float) * 2); - wgpuRenderPassEncoderSetIndexBuffer(renderPassEncoder, context.indexBufferFan, WGPUIndexFormat_Uint32, 0, indexCount * sizeof(uint32_t)); + wgpuRenderPassEncoderSetIndexBuffer(renderPassEncoder, context.bufferIndexFan, WGPUIndexFormat_Uint32, 0, indexCount * sizeof(uint32_t)); wgpuRenderPassEncoderDrawIndexed(renderPassEncoder, indexCount, 1, 0, 0, 0); } @@ -62,8 +60,8 @@ void WgMeshData::update(WgContext& context, const WgPolyline* polyline) assert(polyline->pts.count > 2); vertexCount = polyline->pts.count; indexCount = (polyline->pts.count - 2) * 3; - context.allocateVertexBuffer(bufferPosition, &polyline->pts[0], vertexCount * sizeof(float) * 2); - context.allocateIndexBufferFan(vertexCount); + context.allocateBufferVertex(bufferPosition, (float *)&polyline->pts[0], vertexCount * sizeof(float) * 2); + context.allocateBufferIndexFan(vertexCount); } @@ -74,13 +72,13 @@ void WgMeshData::update(WgContext& context, const WgGeometryData* geometryData) indexCount = geometryData->indexes.count; // buffer position data create and write if (geometryData->positions.pts.count > 0) - context.allocateVertexBuffer(bufferPosition, &geometryData->positions.pts[0], vertexCount * sizeof(float) * 2); + context.allocateBufferVertex(bufferPosition, (float *)&geometryData->positions.pts[0], vertexCount * sizeof(float) * 2); // buffer tex coords data create and write if (geometryData->texCoords.count > 0) - context.allocateVertexBuffer(bufferTexCoord, &geometryData->texCoords[0], vertexCount * sizeof(float) * 2); + context.allocateBufferVertex(bufferTexCoord, (float *)&geometryData->texCoords[0], vertexCount * sizeof(float) * 2); // buffer index data create and write if (geometryData->indexes.count > 0) - context.allocateIndexBuffer(bufferIndex, &geometryData->indexes[0], indexCount * sizeof(uint32_t)); + context.allocateBufferIndex(bufferIndex, &geometryData->indexes[0], indexCount * sizeof(uint32_t)); }; @@ -92,8 +90,8 @@ void WgMeshData::update(WgContext& context, const WgPoint pmin, const WgPoint pm pmin.x, pmin.y, pmax.x, pmin.y, pmax.x, pmax.y, pmin.x, pmax.y }; - context.allocateVertexBuffer(bufferPosition, data, sizeof(data)); - context.allocateIndexBufferFan(vertexCount); + context.allocateBufferVertex(bufferPosition, data, sizeof(data)); + context.allocateBufferIndexFan(vertexCount); } @@ -186,29 +184,14 @@ void WgImageData::update(WgContext& context, Surface* surface) { release(context); assert(surface); - texture = context.createTexture2d( - WGPUTextureUsage_TextureBinding | WGPUTextureUsage_CopyDst, - WGPUTextureFormat_RGBA8Unorm, - surface->w, surface->h, "The shape texture"); + texture = context.createTexture(surface->w, surface->h, WGPUTextureFormat_RGBA8Unorm); assert(texture); - textureView = context.createTextureView2d(texture, "The shape texture view"); + textureView = context.createTextureView(texture); assert(textureView); // update texture data - WGPUImageCopyTexture imageCopyTexture{}; - imageCopyTexture.nextInChain = nullptr; - imageCopyTexture.texture = texture; - imageCopyTexture.mipLevel = 0; - imageCopyTexture.origin = { 0, 0, 0 }; - imageCopyTexture.aspect = WGPUTextureAspect_All; - WGPUTextureDataLayout textureDataLayout{}; - textureDataLayout.nextInChain = nullptr; - textureDataLayout.offset = 0; - textureDataLayout.bytesPerRow = 4 * surface->w; - textureDataLayout.rowsPerImage = surface->h; - WGPUExtent3D writeSize{}; - writeSize.width = surface->w; - writeSize.height = surface->h; - writeSize.depthOrArrayLayers = 1; + WGPUImageCopyTexture imageCopyTexture{ .texture = texture }; + WGPUTextureDataLayout textureDataLayout{ .bytesPerRow = 4 * surface->w, .rowsPerImage = surface->h }; + WGPUExtent3D writeSize{ .width = surface->w, .height = surface->h, .depthOrArrayLayers = 1 }; wgpuQueueWriteTexture(context.queue, &imageCopyTexture, surface->data, 4 * surface->w * surface->h, &textureDataLayout, &writeSize); wgpuQueueSubmit(context.queue, 0, nullptr); }; @@ -228,20 +211,31 @@ void WgRenderSettings::update(WgContext& context, const Fill* fill, const uint8_ { // setup fill properties if ((flags & (RenderUpdateFlag::Gradient)) && fill) { + rasterType = WgRenderRasterType::Gradient; // setup linear fill properties if (fill->type() == Type::LinearGradient) { WgShaderTypeLinearGradient linearGradient((LinearGradient*)fill); - bindGroupLinear.initialize(context.device, context.queue, linearGradient); + if (context.allocateBufferUniform(bufferGroupLinear, &linearGradient, sizeof(linearGradient))) { + context.pipelines->layouts.releaseBindGroup(bindGroupLinear); + bindGroupLinear = context.pipelines->layouts.createBindGroupBuffer1Un(bufferGroupLinear); + } fillType = WgRenderSettingsType::Linear; } else if (fill->type() == Type::RadialGradient) { WgShaderTypeRadialGradient radialGradient((RadialGradient*)fill); - bindGroupRadial.initialize(context.device, context.queue, radialGradient); + if (context.allocateBufferUniform(bufferGroupRadial, &radialGradient, sizeof(radialGradient))) { + context.pipelines->layouts.releaseBindGroup(bindGroupRadial); + bindGroupRadial = context.pipelines->layouts.createBindGroupBuffer1Un(bufferGroupRadial); + } fillType = WgRenderSettingsType::Radial; } skip = false; } else if ((flags & (RenderUpdateFlag::Color)) && !fill) { + rasterType = WgRenderRasterType::Solid; WgShaderTypeSolidColor solidColor(color); - bindGroupSolid.initialize(context.device, context.queue, solidColor); + if (context.allocateBufferUniform(bufferGroupSolid, &solidColor, sizeof(solidColor))) { + context.pipelines->layouts.releaseBindGroup(bindGroupSolid); + bindGroupSolid = context.pipelines->layouts.createBindGroupBuffer1Un(bufferGroupSolid); + } fillType = WgRenderSettingsType::Solid; skip = (color[3] == 0); } @@ -250,9 +244,13 @@ void WgRenderSettings::update(WgContext& context, const Fill* fill, const uint8_ void WgRenderSettings::release(WgContext& context) { - bindGroupSolid.release(); - bindGroupLinear.release(); - bindGroupRadial.release(); + context.pipelines->layouts.releaseBindGroup(bindGroupSolid); + context.pipelines->layouts.releaseBindGroup(bindGroupLinear); + context.pipelines->layouts.releaseBindGroup(bindGroupRadial); + context.releaseBuffer(bufferGroupSolid); + context.releaseBuffer(bufferGroupLinear); + context.releaseBuffer(bufferGroupRadial); + }; //*********************************************************************** @@ -261,11 +259,26 @@ void WgRenderSettings::release(WgContext& context) void WgRenderDataPaint::release(WgContext& context) { - bindGroupPaint.release(); + context.pipelines->layouts.releaseBindGroup(bindGroupPaint); + context.releaseBuffer(bufferModelMat); + context.releaseBuffer(bufferBlendSettings); clips.clear(); }; +void WgRenderDataPaint::update(WgContext& context, const tvg::Matrix& transform, tvg::ColorSpace cs, uint8_t opacity) +{ + WgShaderTypeMat4x4f modelMat(transform); + WgShaderTypeBlendSettings blendSettings(cs, opacity); + bool bufferModelMatChanged = context.allocateBufferUniform(bufferModelMat, &modelMat, sizeof(modelMat)); + bool bufferBlendSettingsChanged = context.allocateBufferUniform(bufferBlendSettings, &blendSettings, sizeof(blendSettings)); + if (bufferModelMatChanged || bufferBlendSettingsChanged) { + context.pipelines->layouts.releaseBindGroup(bindGroupPaint); + bindGroupPaint = context.pipelines->layouts.createBindGroupBuffer2Un(bufferModelMat, bufferBlendSettings); + } +} + + void WgRenderDataPaint::updateClips(tvg::Array &clips) { this->clips.clear(); for (uint32_t i = 0; i < clips.count; i++) @@ -496,8 +509,6 @@ void WgRenderDataPicture::release(WgContext& context) { meshData.release(context); imageData.release(context); - bindGroupPicture.release(); + context.pipelines->layouts.releaseBindGroup(bindGroupPicture); WgRenderDataPaint::release(context); } - -#endif //_TVG_WG_RENDER_DATA_H_ diff --git a/src/renderer/wg_engine/tvgWgRenderData.h b/src/renderer/wg_engine/tvgWgRenderData.h old mode 100644 new mode 100755 index 94742205..48ba0b87 --- a/src/renderer/wg_engine/tvgWgRenderData.h +++ b/src/renderer/wg_engine/tvgWgRenderData.h @@ -20,6 +20,9 @@ * SOFTWARE. */ +#ifndef _TVG_WG_RENDER_DATA_H_ +#define _TVG_WG_RENDER_DATA_H_ + #include "tvgWgPipelines.h" #include "tvgWgGeometry.h" @@ -70,13 +73,18 @@ struct WgImageData { }; enum class WgRenderSettingsType { None = 0, Solid = 1, Linear = 2, Radial = 3 }; +enum class WgRenderRasterType { Solid = 0, Gradient, Image }; struct WgRenderSettings { - WgBindGroupSolidColor bindGroupSolid{}; - WgBindGroupLinearGradient bindGroupLinear{}; - WgBindGroupRadialGradient bindGroupRadial{}; + WGPUBuffer bufferGroupSolid{}; + WGPUBuffer bufferGroupLinear{}; + WGPUBuffer bufferGroupRadial{}; + WGPUBindGroup bindGroupSolid{}; + WGPUBindGroup bindGroupLinear{}; + WGPUBindGroup bindGroupRadial{}; WgRenderSettingsType fillType{}; + WgRenderRasterType rasterType{}; bool skip{}; void update(WgContext& context, const Fill* fill, const uint8_t* color, const RenderUpdateFlag flags); @@ -85,7 +93,9 @@ struct WgRenderSettings struct WgRenderDataPaint { - WgBindGroupPaint bindGroupPaint{}; + WGPUBuffer bufferModelMat{}; + WGPUBuffer bufferBlendSettings{}; + WGPUBindGroup bindGroupPaint{}; RenderRegion viewport{}; float opacity{}; Array clips; @@ -94,6 +104,7 @@ struct WgRenderDataPaint virtual void release(WgContext& context); virtual Type type() { return Type::Undefined; }; + void update(WgContext& context, const tvg::Matrix& transform, tvg::ColorSpace cs, uint8_t opacity); void updateClips(tvg::Array &clips); }; @@ -133,11 +144,12 @@ public: struct WgRenderDataPicture: public WgRenderDataPaint { - WgBindGroupPicture bindGroupPicture{}; + WGPUBindGroup bindGroupPicture{}; WgImageData imageData{}; WgMeshData meshData{}; - void update(WgContext& context); void release(WgContext& context) override; Type type() override { return Type::Picture; }; }; + +#endif // _TVG_WG_RENDER_DATA_H_ diff --git a/src/renderer/wg_engine/tvgWgRenderTarget.cpp b/src/renderer/wg_engine/tvgWgRenderTarget.cpp old mode 100644 new mode 100755 index 09cc5ea4..0e58dd43 --- a/src/renderer/wg_engine/tvgWgRenderTarget.cpp +++ b/src/renderer/wg_engine/tvgWgRenderTarget.cpp @@ -22,407 +22,69 @@ #include "tvgWgRenderTarget.h" -//***************************************************************************** -// render storage -//***************************************************************************** - - void WgRenderStorage::initialize(WgContext& context, uint32_t w, uint32_t h, uint32_t samples, WGPUTextureFormat format) - { - release(context); - // store target storage size - this->samples = samples; - width = w * samples; - height = h * samples; - 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 - texColor = context.createTexture2d( - WGPUTextureUsage_CopySrc | - WGPUTextureUsage_CopyDst | - WGPUTextureUsage_TextureBinding | - WGPUTextureUsage_StorageBinding | - WGPUTextureUsage_RenderAttachment, - format, width, height, "The target texture color"); - texStencil = context.createTexture2d( - WGPUTextureUsage_RenderAttachment, - WGPUTextureFormat_Stencil8, - width, height, "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 - if (format == WGPUTextureFormat_RGBA8Unorm) { - bindGroupTexStorageRgbaRO.initialize(context.device, context.queue, texViewColor); - bindGroupTexStorageRgbaWO.initialize(context.device, context.queue, texViewColor); - } - if (format == WGPUTextureFormat_BGRA8Unorm) - bindGroupTexStorageBgraWO.initialize(context.device, context.queue, texViewColor); - // initialize window binding groups - WgShaderTypeMat4x4f viewMat(w, h); - mBindGroupCanvas.initialize(context.device, context.queue, viewMat); - mPipelines = context.pipelines; - } +void WgRenderStorage::initialize(WgContext& context, uint32_t width, uint32_t height) +{ + this->width = width; + this->height = height; + texture = context.createTexStorage(width, height, WGPUTextureFormat_RGBA8Unorm); + texView = context.createTextureView(texture); + bindGroupRead = context.pipelines->layouts.createBindGroupStrorage1RO(texView); + bindGroupWrite = context.pipelines->layouts.createBindGroupStrorage1WO(texView); + bindGroupTexure = context.pipelines->layouts.createBindGroupTexSampled(context.samplerNearest, texView); +} void WgRenderStorage::release(WgContext& context) { - mRenderPassEncoder = nullptr; - mBindGroupCanvas.release(); - bindGroupTexStorageBgraWO.release(); - bindGroupTexStorageRgbaWO.release(); - bindGroupTexStorageRgbaRO.release(); - context.releaseTextureView(texViewStencil); - context.releaseTextureView(texViewColor); - context.releaseTexture(texStencil); - context.releaseTexture(texColor); - workgroupsCountX = 0; - workgroupsCountY = 0; + context.pipelines->layouts.releaseBindGroup(bindGroupTexure); + context.pipelines->layouts.releaseBindGroup(bindGroupWrite); + context.pipelines->layouts.releaseBindGroup(bindGroupRead); + context.releaseTextureView(texView); + context.releaseTexture(texture); height = 0; width = 0; } - -void WgRenderStorage::renderShape(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType) -{ - assert(renderData); - assert(mRenderPassEncoder); - if (renderData->strokeFirst) { - drawStroke(context, renderData, blendType); - drawShape(context, renderData, blendType); - } else { - drawShape(context, renderData, blendType); - drawStroke(context, renderData, blendType); - } -} - - -void WgRenderStorage::renderPicture(WgContext& context, WgRenderDataPicture* renderData, WgPipelineBlendType blendType) -{ - assert(renderData); - assert(mRenderPassEncoder); - uint8_t blend = (uint8_t)blendType; - auto& vp = renderData->viewport; - if ((vp.w <= 0) || (vp.h <= 0)) return; - wgpuRenderPassEncoderSetScissorRect(mRenderPassEncoder, vp.x * samples, vp.y * samples, vp.w * samples, vp.h * samples); - wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 0); - mPipelines->image[blend].use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint, renderData->bindGroupPicture); - renderData->meshData.drawImage(context, mRenderPassEncoder); -} - -void WgRenderStorage::renderClipPath(WgContext& context, WgRenderDataPaint* renderData) { - assert(renderData); - assert(mRenderPassEncoder); - if (renderData->type() == Type::Shape) - drawShapeClipPath(context, (WgRenderDataShape*)renderData); - else if (renderData->type() == Type::Picture) - drawPictureClipPath(context, (WgRenderDataPicture*)renderData); -} - -void WgRenderStorage::drawShape(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType) -{ - assert(renderData); - assert(mRenderPassEncoder); - assert(renderData->meshGroupShapes.meshes.count == renderData->meshGroupShapesBBox.meshes.count); - if (renderData->renderSettingsShape.skip) return; - // apply viewport - auto& vp = renderData->viewport; - if ((vp.w <= 0) || (vp.h <= 0)) return; - wgpuRenderPassEncoderSetScissorRect(mRenderPassEncoder, vp.x * samples, vp.y * samples, vp.w * samples, vp.h * samples); - // setup fill rule - wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 0); - if (renderData->fillRule == FillRule::Winding) - mPipelines->fillShapeWinding.use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint); - else - mPipelines->fillShapeEvenOdd.use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint); - // draw to stencil (first pass) - for (uint32_t i = 0; i < renderData->meshGroupShapes.meshes.count; i++) - renderData->meshGroupShapes.meshes[i]->drawFan(context, mRenderPassEncoder); - // fill shape geometry (second pass) - uint8_t blend = (uint8_t)blendType; - WgRenderSettings& settings = renderData->renderSettingsShape; - if (settings.fillType == WgRenderSettingsType::Solid) - mPipelines->solid[blend].use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint, settings.bindGroupSolid); - else if (settings.fillType == WgRenderSettingsType::Linear) - mPipelines->linear[blend].use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint, settings.bindGroupLinear); - else if (settings.fillType == WgRenderSettingsType::Radial) - mPipelines->radial[blend].use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint, settings.bindGroupRadial); - renderData->meshDataBBox.drawFan(context, mRenderPassEncoder); -} - - -void WgRenderStorage::drawStroke(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType) -{ - assert(renderData); - assert(mRenderPassEncoder); - assert(renderData->meshGroupStrokes.meshes.count == renderData->meshGroupStrokesBBox.meshes.count); - if (renderData->renderSettingsStroke.skip) return; - // apply viewport - auto& vp = renderData->viewport; - if ((vp.w <= 0) || (vp.h <= 0)) return; - wgpuRenderPassEncoderSetScissorRect(mRenderPassEncoder, vp.x * samples, vp.y * samples, vp.w * samples, vp.h * samples); - // draw stroke geometry - uint8_t blend = (uint8_t)blendType; - // draw strokes to stencil (first pass) - for (uint32_t i = 0; i < renderData->meshGroupStrokes.meshes.count; i++) { - // draw to stencil (first pass) - wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 255); - mPipelines->fillStroke.use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint); - renderData->meshGroupStrokes.meshes[i]->draw(context, mRenderPassEncoder); - // fill shape (second pass) - wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 0); - WgRenderSettings& settings = renderData->renderSettingsStroke; - if (settings.fillType == WgRenderSettingsType::Solid) - mPipelines->solid[blend].use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint, settings.bindGroupSolid); - else if (settings.fillType == WgRenderSettingsType::Linear) - mPipelines->linear[blend].use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint, settings.bindGroupLinear); - else if (settings.fillType == WgRenderSettingsType::Radial) - mPipelines->radial[blend].use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint, settings.bindGroupRadial); - renderData->meshGroupStrokesBBox.meshes[i]->drawFan(context, mRenderPassEncoder); - } -} - - -void WgRenderStorage::drawShapeClipPath(WgContext& context, WgRenderDataShape* renderData) { - assert(renderData); - assert(renderData->type() == Type::Shape); - assert(renderData->meshGroupShapes.meshes.count == renderData->meshGroupShapesBBox.meshes.count); - // draw shape geometry - wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 0); - // setup fill rule - if (renderData->fillRule == FillRule::Winding) - mPipelines->fillShapeWinding.use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint); - else - mPipelines->fillShapeEvenOdd.use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint); - // draw to stencil (first pass) - for (uint32_t i = 0; i < renderData->meshGroupShapes.meshes.count; i++) - renderData->meshGroupShapes.meshes[i]->drawFan(context, mRenderPassEncoder); - // fill shape geometry (second pass) - mPipelines->clipMask.use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint); - renderData->meshDataBBox.drawFan(context, mRenderPassEncoder); -} - -void WgRenderStorage::drawPictureClipPath(WgContext& context, WgRenderDataPicture* renderData) { - assert(renderData); - assert(renderData->type() == Type::Picture); - assert(mRenderPassEncoder); - wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 0); - mPipelines->clipMask.use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint); - renderData->meshData.drawImage(context, mRenderPassEncoder); -} - - - -void WgRenderStorage::blend( - WgContext& context, - WGPUCommandEncoder commandEncoder, - WgPipelineBlend* pipeline, - WgRenderStorage* texSrc, - WgRenderStorage* texDst, - WgBindGroupBlendMethod* blendMethod, - WgBindGroupOpacity* opacity) -{ - assert(commandEncoder); - assert(texSrc); - assert(texDst); - assert(pipeline); - WgBindGroupTexBlend texBlend; - texBlend.initialize(context.device, context.queue, texSrc->texViewColor, texDst->texViewColor, texViewColor); - WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder); - pipeline->use(computePassEncoder, texBlend, *blendMethod, *opacity); - dispatchWorkgroups(computePassEncoder); - endComputePass(computePassEncoder); -} - - -void WgRenderStorage::blendMask( - WgContext& context, - WGPUCommandEncoder commandEncoder, - WgPipelineBlendMask* pipeline, - WgRenderStorage* texMsk, - WgRenderStorage* texSrc, - WgRenderStorage* texDst, - WgBindGroupBlendMethod* blendMethod, - WgBindGroupOpacity* opacity) -{ - assert(commandEncoder); - assert(texSrc); - assert(texMsk); - WgBindGroupTexBlendMask texBlendMask; - texBlendMask.initialize(context.device, context.queue, texSrc->texViewColor, texMsk->texViewColor, texDst->texViewColor, texViewColor); - WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder); - pipeline->use(computePassEncoder, texBlendMask, *blendMethod, *opacity); - dispatchWorkgroups(computePassEncoder); - endComputePass(computePassEncoder); - texBlendMask.release(); -}; - - -void WgRenderStorage::maskCompose( - WgContext& context, - WGPUCommandEncoder commandEncoder, - WgRenderStorage* texMsk0, - WgRenderStorage* texMsk1) -{ - assert(commandEncoder); - assert(texMsk0); - WgBindGroupTexMaskCompose maskCompose; - maskCompose.initialize(context.device, context.queue, texMsk0->texViewColor, texMsk1->texViewColor, texViewColor); - WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder); - mPipelines->computeMaskCompose.use(computePassEncoder, maskCompose); - dispatchWorkgroups(computePassEncoder); - endComputePass(computePassEncoder); - maskCompose.release(); -} - - -void WgRenderStorage::compose( - WgContext& context, - WGPUCommandEncoder commandEncoder, - WgRenderStorage* texSrc, - WgRenderStorage* texMsk, - WgRenderStorage* texDst, - WgBindGroupCompositeMethod* composeMethod, - WgBindGroupBlendMethod* blendMethod, - WgBindGroupOpacity* opacity) -{ - assert(commandEncoder); - assert(texSrc); - assert(texMsk); - WgBindGroupTexCompose texCompose; - texCompose.initialize(context.device, context.queue, texSrc->texViewColor, texMsk->texViewColor, texDst->texViewColor, texViewColor); - WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder); - mPipelines->computeCompose.use(computePassEncoder, texCompose, *composeMethod, *blendMethod, *opacity); - dispatchWorkgroups(computePassEncoder); - endComputePass(computePassEncoder); - texCompose.release(); -} - - -void WgRenderStorage::antialias(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc) -{ - assert(commandEncoder); - assert(targetSrc); - WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder); - mPipelines->computeAntiAliasing.use(computePassEncoder, targetSrc->bindGroupTexStorageRgbaRO, bindGroupTexStorageBgraWO); - dispatchWorkgroups(computePassEncoder); - endComputePass(computePassEncoder); -} - -void WgRenderStorage::copy(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc) -{ - assert(commandEncoder); - assert(targetSrc); - WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder); - mPipelines->computeCopy.use(computePassEncoder, targetSrc->bindGroupTexStorageRgbaRO, bindGroupTexStorageRgbaWO); - dispatchWorkgroups(computePassEncoder); - endComputePass(computePassEncoder); -} - -void WgRenderStorage::dispatchWorkgroups(WGPUComputePassEncoder computePassEncoder) -{ - assert(computePassEncoder); - wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, workgroupsCountX, workgroupsCountY, 1); -} - - -void WgRenderStorage::beginRenderPass(WGPUCommandEncoder commandEncoder, bool clear) -{ - assert(commandEncoder); - // render pass depth stencil attachment - WGPURenderPassDepthStencilAttachment depthStencilAttachment{}; - depthStencilAttachment.view = texViewStencil; - depthStencilAttachment.depthLoadOp = WGPULoadOp_Undefined; - depthStencilAttachment.depthStoreOp = WGPUStoreOp_Undefined; - depthStencilAttachment.depthClearValue = 1.0f; - depthStencilAttachment.depthReadOnly = false; - depthStencilAttachment.stencilLoadOp = WGPULoadOp_Clear; - depthStencilAttachment.stencilStoreOp = WGPUStoreOp_Discard; - depthStencilAttachment.stencilClearValue = 0; - depthStencilAttachment.stencilReadOnly = false; - // render pass color attachment - WGPURenderPassColorAttachment colorAttachment{}; - colorAttachment.view = texViewColor; - #ifdef __EMSCRIPTEN__ - colorAttachment.depthSlice = WGPU_DEPTH_SLICE_UNDEFINED; - #endif - colorAttachment.resolveTarget = nullptr; - colorAttachment.loadOp = clear ? WGPULoadOp_Clear : WGPULoadOp_Load; - colorAttachment.clearValue = { 0, 0, 0, 0 }; - colorAttachment.storeOp = WGPUStoreOp_Store; - // render pass descriptor - WGPURenderPassDescriptor renderPassDesc{}; - renderPassDesc.nextInChain = nullptr; - renderPassDesc.label = "The render pass"; - renderPassDesc.colorAttachmentCount = 1; - renderPassDesc.colorAttachments = &colorAttachment; - renderPassDesc.depthStencilAttachment = &depthStencilAttachment; - //renderPassDesc.depthStencilAttachment = nullptr; - renderPassDesc.occlusionQuerySet = nullptr; - renderPassDesc.timestampWrites = nullptr; - // begin render pass - mRenderPassEncoder = wgpuCommandEncoderBeginRenderPass(commandEncoder, &renderPassDesc); -} - - -void WgRenderStorage::endRenderPass() -{ - assert(mRenderPassEncoder); - wgpuRenderPassEncoderEnd(mRenderPassEncoder); - wgpuRenderPassEncoderRelease(mRenderPassEncoder); - mRenderPassEncoder = nullptr; -} - - -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::endComputePass(WGPUComputePassEncoder computePassEncoder) -{ - assert(computePassEncoder); - wgpuComputePassEncoderEnd(computePassEncoder); -} - //***************************************************************************** // render storage pool //***************************************************************************** -WgRenderStorage* WgRenderStoragePool::allocate(WgContext& context, uint32_t w, uint32_t h, uint32_t samples) +WgRenderStorage* WgRenderStoragePool::allocate(WgContext& context) { - WgRenderStorage* renderStorage{}; - if (mPool.count > 0) { - renderStorage = mPool.last(); - mPool.pop(); - } else { - renderStorage = new WgRenderStorage; - renderStorage->initialize(context, w, h, samples); - mList.push(renderStorage); - } - return renderStorage; + WgRenderStorage* renderStorage{}; + if (pool.count > 0) { + renderStorage = pool.last(); + pool.pop(); + } else { + renderStorage = new WgRenderStorage; + renderStorage->initialize(context, width, height); + list.push(renderStorage); + } + return renderStorage; }; -void WgRenderStoragePool::free(WgContext& context, WgRenderStorage* renderStorage) { - mPool.push(renderStorage); +void WgRenderStoragePool::free(WgContext& context, WgRenderStorage* renderStorage) +{ + pool.push(renderStorage); }; +void WgRenderStoragePool::initialize(WgContext& context, uint32_t width, uint32_t height) +{ + this->width = width; + this->height = height; +} + + void WgRenderStoragePool::release(WgContext& context) { - for (uint32_t i = 0; i < mList.count; i++) { - mList[i]->release(context); - delete mList[i]; - } - mList.clear(); - mPool.clear(); + for (uint32_t i = 0; i < list.count; i++) { + list[i]->release(context); + delete list[i]; + } + list.clear(); + pool.clear(); + height = 0; + width = 0; }; diff --git a/src/renderer/wg_engine/tvgWgRenderTarget.h b/src/renderer/wg_engine/tvgWgRenderTarget.h old mode 100644 new mode 100755 index 70e4b87f..332d616b --- a/src/renderer/wg_engine/tvgWgRenderTarget.h +++ b/src/renderer/wg_engine/tvgWgRenderTarget.h @@ -23,94 +23,34 @@ #ifndef _TVG_WG_RENDER_TARGET_H_ #define _TVG_WG_RENDER_TARGET_H_ -#include "tvgWgRenderData.h" +#include "tvgWgPipelines.h" +#include "tvgRender.h" -class WgRenderStorage { -private: - // texture buffers - WgBindGroupCanvas mBindGroupCanvas; - WGPURenderPassEncoder mRenderPassEncoder{}; - WgPipelines* mPipelines{}; // external handle -public: - WGPUTexture texColor{}; - WGPUTexture texStencil{}; - WGPUTextureView texViewColor{}; - WGPUTextureView texViewStencil{}; - WgBindGroupTextureStorageRgbaRO bindGroupTexStorageRgbaRO; - WgBindGroupTextureStorageRgbaWO bindGroupTexStorageRgbaWO; - WgBindGroupTextureStorageBgraWO bindGroupTexStorageBgraWO; - uint32_t samples{}; +struct WgRenderStorage { + WGPUTexture texture{}; + WGPUTextureView texView{}; + WGPUBindGroup bindGroupRead{}; + WGPUBindGroup bindGroupWrite{}; + WGPUBindGroup bindGroupTexure{}; uint32_t width{}; uint32_t height{}; - uint32_t workgroupsCountX{}; - uint32_t workgroupsCountY{}; -public: - void initialize(WgContext& context, uint32_t w, uint32_t h, uint32_t samples = 1, WGPUTextureFormat format = WGPUTextureFormat_RGBA8Unorm); + + void initialize(WgContext& context, uint32_t width, uint32_t height); void release(WgContext& context); - - void beginRenderPass(WGPUCommandEncoder commandEncoder, bool clear); - void endRenderPass(); - - void renderShape(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType); - void renderPicture(WgContext& context, WgRenderDataPicture* renderData, WgPipelineBlendType blendType); - void renderClipPath(WgContext& context, WgRenderDataPaint* renderData); - - void blend( - WgContext& context, - WGPUCommandEncoder commandEncoder, - WgPipelineBlend* pipeline, - WgRenderStorage* targetSrc, - WgRenderStorage* targetDst, - WgBindGroupBlendMethod* blendMethod, - WgBindGroupOpacity* opacity); - void blendMask( - WgContext& context, - WGPUCommandEncoder commandEncoder, - WgPipelineBlendMask* pipeline, - WgRenderStorage* texMsk, - WgRenderStorage* texSrc, - WgRenderStorage* texDst, - WgBindGroupBlendMethod* blendMethod, - WgBindGroupOpacity* opacity); - void maskCompose( - WgContext& context, - WGPUCommandEncoder commandEncoder, - WgRenderStorage* texMsk0, - WgRenderStorage* texMsk1); - void compose( - WgContext& context, - WGPUCommandEncoder commandEncoder, - WgRenderStorage* texMsk, - WgRenderStorage* texSrc, - WgRenderStorage* texDst, - WgBindGroupCompositeMethod* composeMethod, - WgBindGroupBlendMethod* blendMethod, - WgBindGroupOpacity* opacity); - void antialias(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc); - void copy(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc); -private: - void drawShape(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType); - void drawStroke(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType); - - void drawShapeClipPath(WgContext& context, WgRenderDataShape* renderData); - void drawPictureClipPath(WgContext& context, WgRenderDataPicture* renderData); - - void dispatchWorkgroups(WGPUComputePassEncoder computePassEncoder); - - WGPUComputePassEncoder beginComputePass(WGPUCommandEncoder commandEncoder); - void endComputePass(WGPUComputePassEncoder computePassEncoder); }; class WgRenderStoragePool { private: - Array mList; - Array mPool; + Array list; + Array pool; + uint32_t width{}; + uint32_t height{}; public: - WgRenderStorage* allocate(WgContext& context, uint32_t w, uint32_t h, uint32_t samples = 1); - void free(WgContext& context, WgRenderStorage* renderTarget); - void release(WgContext& context); + WgRenderStorage* allocate(WgContext& context); + void free(WgContext& context, WgRenderStorage* renderTarget); + + void initialize(WgContext& context, uint32_t width, uint32_t height); + void release(WgContext& context); }; - - -#endif +#endif // _TVG_WG_RENDER_TARGET_H_ diff --git a/src/renderer/wg_engine/tvgWgRenderer.cpp b/src/renderer/wg_engine/tvgWgRenderer.cpp old mode 100644 new mode 100755 index 40961688..70b049f1 --- a/src/renderer/wg_engine/tvgWgRenderer.cpp +++ b/src/renderer/wg_engine/tvgWgRenderer.cpp @@ -21,15 +21,12 @@ */ #include "tvgWgRenderer.h" - -#define WG_SSAA_SAMPLES (2) - -/************************************************************************/ -/* Internal Class Implementation */ -/************************************************************************/ +#include WgRenderer::WgRenderer() { + WgGeometryData::gMath = new WgMath(); + WgGeometryData::gMath->initialize(); } @@ -37,49 +34,40 @@ WgRenderer::~WgRenderer() { release(); mContext.release(); + WgGeometryData::gMath->release(); + delete WgGeometryData::gMath; } void WgRenderer::initialize() { mPipelines.initialize(mContext); - mOpacityPool.initialize(mContext); - mBlendMethodPool.initialize(mContext); - mCompositeMethodPool.initialize(mContext); WgMeshDataGroup::gMeshDataPool = new WgMeshDataPool(); - WgGeometryData::gMath = new WgMath(); - WgGeometryData::gMath->initialize(); } void WgRenderer::release() { - clearDisposes(); - WgGeometryData::gMath->release(); - delete WgGeometryData::gMath; + disposeObjects(); + mContext.releaseTexture(mTexScreen); + mContext.releaseTextureView(mTexViewScreen); + mContext.pipelines->layouts.releaseBindGroup(mBindGroupScreen); + mStorageRoot.release(mContext); + mRenderStoragePool.release(mContext); mRenderDataShapePool.release(mContext); WgMeshDataGroup::gMeshDataPool->release(mContext); delete WgMeshDataGroup::gMeshDataPool; mCompositorStack.clear(); mRenderStorageStack.clear(); - mRenderStoragePool.release(mContext); - mCompositeMethodPool.release(mContext); - mBlendMethodPool.release(mContext); - mOpacityPool.release(mContext); - mRenderStorageRoot.release(mContext); - mRenderStorageMask.release(mContext); - mRenderStorageCopy.release(mContext); - mRenderStorageScreen.release(mContext); - mRenderStorageInterm.release(mContext); - mPipelines.release(); + mPipelines.release(mContext); } -void WgRenderer::clearDisposes() +void WgRenderer::disposeObjects() { - if (mDisposed.renderDatas.count == 0) return; + if (mDisposeRenderDatas.count == 0) return; - for (auto p = mDisposed.renderDatas.begin(); p < mDisposed.renderDatas.end(); p++) { + for (auto p = mDisposeRenderDatas.begin(); p < mDisposeRenderDatas.end(); p++) { auto renderData = (WgRenderDataPaint*)(*p); if (renderData->type() == Type::Shape) { mRenderDataShapePool.free(mContext, (WgRenderDataShape*)renderData); @@ -87,12 +75,9 @@ void WgRenderer::clearDisposes() renderData->release(mContext); } } - mDisposed.renderDatas.clear(); + mDisposeRenderDatas.clear(); } -/************************************************************************/ -/* External Class Implementation */ -/************************************************************************/ RenderData WgRenderer::prepare(const RenderShape& rshape, RenderData data, const Matrix& transform, Array& clips, uint8_t opacity, RenderUpdateFlag flags, bool clipper) { @@ -106,11 +91,9 @@ RenderData WgRenderer::prepare(const RenderShape& rshape, RenderData data, const renderDataShape->updateMeshes(mContext, rshape, transform); } - // update paint settings + // update paint settings if ((!data) || (flags & (RenderUpdateFlag::Transform | RenderUpdateFlag::Blend))) { - WgShaderTypeMat4x4f modelMat(transform); - WgShaderTypeBlendSettings blendSettings(mTargetSurface.cs, opacity); - renderDataShape->bindGroupPaint.initialize(mContext.device, mContext.queue, modelMat, blendSettings); + renderDataShape->update(mContext, transform, mTargetSurface.cs, opacity); renderDataShape->fillRule = rshape.rule; } @@ -145,9 +128,7 @@ RenderData WgRenderer::prepare(Surface* surface, const RenderMesh* mesh, RenderD renderDataPicture->viewport = mViewport; renderDataPicture->opacity = opacity; if (flags & (RenderUpdateFlag::Transform | RenderUpdateFlag::Blend)) { - WgShaderTypeMat4x4f modelMat(transform); - WgShaderTypeBlendSettings blendSettings(surface->cs, opacity); - renderDataPicture->bindGroupPaint.initialize(mContext.device, mContext.queue, modelMat, blendSettings); + renderDataPicture->update(mContext, transform, surface->cs, opacity); } // update image data @@ -155,13 +136,13 @@ RenderData WgRenderer::prepare(Surface* surface, const RenderMesh* mesh, RenderD WgGeometryData geometryData; if (mesh->triangleCnt == 0) geometryData.appendImageBox(surface->w, surface->h); else geometryData.appendMesh(mesh); + mContext.pipelines->layouts.releaseBindGroup(renderDataPicture->bindGroupPicture); renderDataPicture->meshData.release(mContext); renderDataPicture->meshData.update(mContext, &geometryData); renderDataPicture->imageData.update(mContext, surface); - renderDataPicture->bindGroupPicture.initialize( - mContext.device, mContext.queue, - mContext.samplerLinear, - renderDataPicture->imageData.textureView); + renderDataPicture->bindGroupPicture = mContext.pipelines->layouts.createBindGroupTexSampled( + mContext.samplerLinear, renderDataPicture->imageData.textureView + ); } // store clips data @@ -172,151 +153,47 @@ RenderData WgRenderer::prepare(Surface* surface, const RenderMesh* mesh, RenderD bool WgRenderer::preRender() { + // push rot render storage to the render tree stack + assert(mRenderStorageStack.count == 0); + mRenderStorageStack.push(&mStorageRoot); + // create command encoder for drawing WGPUCommandEncoderDescriptor commandEncoderDesc{}; - commandEncoderDesc.nextInChain = nullptr; - commandEncoderDesc.label = "The command encoder"; mCommandEncoder = wgpuDeviceCreateCommandEncoder(mContext.device, &commandEncoderDesc); - mRenderStorageStack.push(&mRenderStorageRoot); - //mRenderStorageRoot.clear(mCommandEncoder); - mRenderStorageRoot.beginRenderPass(mCommandEncoder, true); + // start root render pass + mCompositor.beginRenderPass(mCommandEncoder, mRenderStorageStack.last(), true); return true; } -void WgRenderer::renderClipPath(Array& clips) -{ - for (uint32_t i = 0; i < clips.count; i++) { - renderClipPath(clips[i]->clips); - // render image to render target - mRenderStorageInterm.beginRenderPass(mCommandEncoder, true); - mRenderStorageInterm.renderClipPath(mContext, clips[i]); - mRenderStorageInterm.endRenderPass(); - mRenderStorageCopy.copy(mCommandEncoder, &mRenderStorageMask); - mRenderStorageMask.maskCompose(mContext, mCommandEncoder, &mRenderStorageInterm, &mRenderStorageCopy); - } -} - - bool WgRenderer::renderShape(RenderData data) { - // get current render storage - WgRenderDataShape *dataShape = (WgRenderDataShape*)data; - if (dataShape->opacity == 0) return 0; - WgPipelineBlendType blendType = WgPipelines::blendMethodToBlendType(mBlendMethod); - WgRenderStorage* renderStorage = mRenderStorageStack.last(); - assert(renderStorage); - // use masked blend - if (dataShape->clips.count > 0) { - // terminate current render pass - renderStorage->endRenderPass(); - // render clip path - mRenderStorageMask.beginRenderPass(mCommandEncoder, true); - mRenderStorageMask.renderClipPath(mContext, dataShape->clips[0]); - mRenderStorageMask.endRenderPass(); - renderClipPath(dataShape->clips); - // render image to render target - mRenderStorageInterm.beginRenderPass(mCommandEncoder, true); - mRenderStorageInterm.renderShape(mContext, dataShape, blendType); - mRenderStorageInterm.endRenderPass(); - // blend shape with current render storage - WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod); - WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, 255); - WgPipelineBlendMask* pipeline = &mContext.pipelines->computeBlendSolidMask; - if (dataShape->renderSettingsShape.fillType != WgRenderSettingsType::Solid) - pipeline = &mContext.pipelines->computeBlendGradientMask; - mRenderStorageCopy.copy(mCommandEncoder, renderStorage); - renderStorage->blendMask(mContext, mCommandEncoder, - pipeline, &mRenderStorageMask, &mRenderStorageInterm, &mRenderStorageCopy, blendMethod, opacity); - // restore current render pass - renderStorage->beginRenderPass(mCommandEncoder, false); - // use hardware blend - } else if (WgPipelines::isBlendMethodSupportsHW(mBlendMethod)) { - renderStorage->renderShape(mContext, dataShape, blendType); - // use custom blend - } else { - // terminate current render pass - renderStorage->endRenderPass(); - // render image to render target - mRenderStorageInterm.beginRenderPass(mCommandEncoder, true); - mRenderStorageInterm.renderShape(mContext, dataShape, blendType); - mRenderStorageInterm.endRenderPass(); - // blend shape with current render storage - WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod); - WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, 255); - WgPipelineBlend* pipeline = &mContext.pipelines->computeBlendSolid; - if (dataShape->renderSettingsShape.fillType != WgRenderSettingsType::Solid) - pipeline = &mContext.pipelines->computeBlendGradient; - mRenderStorageCopy.copy(mCommandEncoder, renderStorage); - renderStorage->blend(mContext, mCommandEncoder, - pipeline, &mRenderStorageInterm, &mRenderStorageCopy, blendMethod, opacity); - // restore current render pass - renderStorage->beginRenderPass(mCommandEncoder, false); - } + // temporary simple render data to the current render target + mCompositor.renderShape(mContext, (WgRenderDataShape*)data, mBlendMethod); return true; } bool WgRenderer::renderImage(RenderData data) { - // get current render storage - WgRenderDataPicture *dataPicture = (WgRenderDataPicture*)data; - // get current render storage - WgPipelineBlendType blendType = WgPipelines::blendMethodToBlendType(mBlendMethod); - WgRenderStorage* renderStorage = mRenderStorageStack.last(); - assert(renderStorage); - // use masked blend - if (dataPicture->clips.count > 0) { - // terminate current render pass - renderStorage->endRenderPass(); - // render clip path - mRenderStorageMask.beginRenderPass(mCommandEncoder, true); - mRenderStorageMask.renderClipPath(mContext, dataPicture->clips[0]); - mRenderStorageMask.endRenderPass(); - renderClipPath(dataPicture->clips); - // render image to render target - mRenderStorageInterm.beginRenderPass(mCommandEncoder, true); - mRenderStorageInterm.renderPicture(mContext, dataPicture, blendType); - mRenderStorageInterm.endRenderPass(); - // blend shape with current render storage - WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod); - WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, 255); - WgPipelineBlendMask* pipeline = &mContext.pipelines->computeBlendImageMask; - mRenderStorageCopy.copy(mCommandEncoder, renderStorage); - renderStorage->blendMask(mContext, mCommandEncoder, - pipeline, &mRenderStorageMask, &mRenderStorageInterm, &mRenderStorageCopy, blendMethod, opacity); - // restore current render pass - renderStorage->beginRenderPass(mCommandEncoder, false); - // use hardware blend - } else if (WgPipelines::isBlendMethodSupportsHW(mBlendMethod)) - renderStorage->renderPicture(mContext, dataPicture, blendType); - // use custom blend - else { - // terminate current render pass - renderStorage->endRenderPass(); - // render image to render target - mRenderStorageInterm.beginRenderPass(mCommandEncoder, true); - mRenderStorageInterm.renderPicture(mContext, dataPicture, blendType); - mRenderStorageInterm.endRenderPass(); - // blend shape with current render storage - WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod); - WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, 255); - WgPipelineBlend* pipeline = &mContext.pipelines->computeBlendImage; - mRenderStorageCopy.copy(mCommandEncoder, renderStorage); - renderStorage->blend(mContext, mCommandEncoder, - pipeline, &mRenderStorageInterm, &mRenderStorageCopy, blendMethod, opacity); - // restore current render pass - renderStorage->beginRenderPass(mCommandEncoder, false); - } + // temporary simple render data to the current render target + mCompositor.renderImage(mContext, (WgRenderDataPicture*)data, mBlendMethod); return true; } bool WgRenderer::postRender() { - mRenderStorageRoot.endRenderPass(); - mRenderStorageStack.pop(); - mContext.executeCommandEncoder(mCommandEncoder); + // end root render pass + mCompositor.endRenderPass(); + // release command encoder + const WGPUCommandBufferDescriptor commandBufferDesc{}; + WGPUCommandBuffer commandsBuffer = wgpuCommandEncoderFinish(mCommandEncoder, &commandBufferDesc); + wgpuQueueSubmit(mContext.queue, 1, &commandsBuffer); + wgpuCommandBufferRelease(commandsBuffer); wgpuCommandEncoderRelease(mCommandEncoder); + // pop root render storage to the render tree stack + mRenderStorageStack.pop(); + assert(mRenderStorageStack.count == 0); return true; } @@ -324,8 +201,8 @@ bool WgRenderer::postRender() void WgRenderer::dispose(RenderData data) { auto renderData = (WgRenderDataPaint*)data; if (renderData) { - ScopedLock lock(mDisposed.key); - mDisposed.renderDatas.push(data); + ScopedLock lock(mDisposeKey); + mDisposeRenderDatas.push(data); } } @@ -375,28 +252,35 @@ bool WgRenderer::clear() bool WgRenderer::sync() { - clearDisposes(); - WGPUSurfaceTexture backBuffer{}; - wgpuSurfaceGetCurrentTexture(mContext.surface, &backBuffer); - - WGPUCommandEncoderDescriptor commandEncoderDesc{}; - commandEncoderDesc.nextInChain = nullptr; - commandEncoderDesc.label = "The command encoder"; + disposeObjects(); + // get current texture + WGPUSurfaceTexture surfaceTexture{}; + wgpuSurfaceGetCurrentTexture(mContext.surface, &surfaceTexture); + + // create command encoder + const WGPUCommandEncoderDescriptor commandEncoderDesc{}; WGPUCommandEncoder commandEncoder = wgpuDeviceCreateCommandEncoder(mContext.device, &commandEncoderDesc); - mRenderStorageScreen.antialias(commandEncoder, &mRenderStorageRoot); + // copy render to screen with conversion rgba to bgra (screen format) + const WGPUComputePassDescriptor computePassDescriptor{}; + WGPUComputePassEncoder computePassEncoder = wgpuCommandEncoderBeginComputePass(commandEncoder, &computePassDescriptor); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 0, mStorageRoot.bindGroupRead, 0, nullptr); + wgpuComputePassEncoderSetBindGroup(computePassEncoder, 1, mBindGroupScreen, 0, nullptr); + wgpuComputePassEncoderSetPipeline(computePassEncoder, mContext.pipelines->copy); + wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, (mStorageRoot.width + 7) / 8, (mStorageRoot.height + 7) / 8, 1); + wgpuComputePassEncoderEnd(computePassEncoder); - WGPUImageCopyTexture source{}; - source.texture = mRenderStorageScreen.texColor; - WGPUImageCopyTexture dest{}; - dest.texture = backBuffer.texture; - WGPUExtent3D copySize{}; - copySize.width = mTargetSurface.w; - copySize.height = mTargetSurface.h; - copySize.depthOrArrayLayers = 1; - wgpuCommandEncoderCopyTextureToTexture(commandEncoder, &source, &dest, ©Size); + // copy dst storage to temporary read only storage + const WGPUImageCopyTexture texSrc { .texture = mTexScreen }; + const WGPUImageCopyTexture texDst { .texture = surfaceTexture.texture }; + const WGPUExtent3D copySize { .width = mTargetSurface.w, .height = mTargetSurface.h, .depthOrArrayLayers = 1 }; + wgpuCommandEncoderCopyTextureToTexture(commandEncoder, &texSrc, &texDst, ©Size); - mContext.executeCommandEncoder(commandEncoder); + // release command encoder + const WGPUCommandBufferDescriptor commandBufferDesc{}; + WGPUCommandBuffer commandsBuffer = wgpuCommandEncoderFinish(commandEncoder, &commandBufferDesc); + wgpuQueueSubmit(mContext.queue, 1, &commandsBuffer); + wgpuCommandBufferRelease(commandsBuffer); wgpuCommandEncoderRelease(commandEncoder); return true; @@ -413,114 +297,88 @@ bool WgRenderer::target(WGPUInstance instance, WGPUSurface surface, uint32_t w, mContext.initialize(instance, surface); - // configure surface - WGPUSurfaceConfiguration surfaceConfiguration{}; - surfaceConfiguration.nextInChain = nullptr; - surfaceConfiguration.device = mContext.device; - surfaceConfiguration.format = WGPUTextureFormat_BGRA8Unorm; - surfaceConfiguration.usage = WGPUTextureUsage_CopyDst; - surfaceConfiguration.viewFormatCount = 0; - surfaceConfiguration.viewFormats = nullptr; - surfaceConfiguration.alphaMode = WGPUCompositeAlphaMode_Auto; - surfaceConfiguration.width = w; - surfaceConfiguration.height = h; - #ifdef __EMSCRIPTEN__ - surfaceConfiguration.presentMode = WGPUPresentMode_Fifo; - #else - surfaceConfiguration.presentMode = WGPUPresentMode_Immediate; - #endif - wgpuSurfaceConfigure(mContext.surface, &surfaceConfiguration); - + WGPUSurfaceConfiguration surfaceConfiguration { + .device = mContext.device, + .format = WGPUTextureFormat_BGRA8Unorm, + .usage = WGPUTextureUsage_CopyDst, + .width = w, .height = h, + #ifdef __EMSCRIPTEN__ + .presentMode = WGPUPresentMode_Fifo, + #else + .presentMode = WGPUPresentMode_Immediate + #endif + }; + wgpuSurfaceConfigure(surface, &surfaceConfiguration); initialize(); - - mRenderStorageInterm.initialize(mContext, w, h, WG_SSAA_SAMPLES); - mRenderStorageMask.initialize(mContext, w, h, WG_SSAA_SAMPLES); - mRenderStorageCopy.initialize(mContext, w, h, WG_SSAA_SAMPLES); - mRenderStorageRoot.initialize(mContext, w, h, WG_SSAA_SAMPLES); - mRenderStorageScreen.initialize(mContext, w, h, 1, WGPUTextureFormat_BGRA8Unorm); + mRenderStoragePool.initialize(mContext, w, h); + mStorageRoot.initialize(mContext, w, h); + mCompositor.initialize(mContext, w, h); + // screen buffer + mTexScreen = mContext.createTexStorage(w, h, WGPUTextureFormat_BGRA8Unorm); + mTexViewScreen = mContext.createTextureView(mTexScreen); + mBindGroupScreen = mContext.pipelines->layouts.createBindGroupScreen1WO(mTexViewScreen); return true; } Compositor* WgRenderer::target(TVG_UNUSED const RenderRegion& region, TVG_UNUSED ColorSpace cs) { - mCompositorStack.push(new WgCompositor); + mCompositorStack.push(new WgCompose); return mCompositorStack.last(); } -bool WgRenderer::beginComposite(TVG_UNUSED Compositor* cmp, TVG_UNUSED CompositeMethod method, TVG_UNUSED uint8_t opacity) +bool WgRenderer::beginComposite(Compositor* cmp, CompositeMethod method, uint8_t opacity) { // save current composition settings - WgCompositor *comp = (WgCompositor*)cmp; - comp->method = method; - comp->opacity = opacity; - comp->blendMethod = mBlendMethod; - + WgCompose* compose = (WgCompose *)cmp; + compose->method = method; + compose->opacity = opacity; + compose->blend = mBlendMethod; // end current render pass - mRenderStorageStack.last()->endRenderPass(); - // allocate new render storage and push it to top of render tree - WgRenderStorage* renderStorage = mRenderStoragePool.allocate(mContext, mTargetSurface.w, mTargetSurface.h, WG_SSAA_SAMPLES); - mRenderStorageStack.push(renderStorage); - // begin last render pass - mRenderStorageStack.last()->beginRenderPass(mCommandEncoder, true); + mCompositor.endRenderPass(); + // allocate new render storage and push to the render tree stack + WgRenderStorage* storage = mRenderStoragePool.allocate(mContext); + mRenderStorageStack.push(storage); + // begin newly added render pass + mCompositor.beginRenderPass(mCommandEncoder, mRenderStorageStack.last(), true); return true; } -bool WgRenderer::endComposite(TVG_UNUSED Compositor* cmp) +bool WgRenderer::endComposite(Compositor* cmp) { - WgCompositor *comp = (WgCompositor*)cmp; + // get current composition settings + WgCompose* comp = (WgCompose *)cmp; + // end current render pass + mCompositor.endRenderPass(); + // finish scene blending if (comp->method == CompositeMethod::None) { - // end current render pass - mRenderStorageStack.last()->endRenderPass(); - // get two last render targets - WgRenderStorage* renderStorageSrc = mRenderStorageStack.last(); + // get source and destination render storages + WgRenderStorage* src = mRenderStorageStack.last(); mRenderStorageStack.pop(); - - // blent scene to current render storage - WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, comp->blendMethod); - WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, comp->opacity); - mRenderStorageCopy.copy(mCommandEncoder, mRenderStorageStack.last()); - mRenderStorageStack.last()->blend(mContext, mCommandEncoder, - &mContext.pipelines->computeBlendImage, - renderStorageSrc, &mRenderStorageCopy, blendMethod, opacity); - + WgRenderStorage* dst = mRenderStorageStack.last(); + // apply blend + mCompositor.blend(mCommandEncoder, src, dst, comp->opacity, comp->blend, WgRenderRasterType::Image); // back render targets to the pool - mRenderStoragePool.free(mContext, renderStorageSrc); - // begin last render pass - mRenderStorageStack.last()->beginRenderPass(mCommandEncoder, false); - } else { - // end current render pass - mRenderStorageStack.last()->endRenderPass(); - + mRenderStoragePool.free(mContext, src); + } else { // finish composition // get source, mask and destination render storages - WgRenderStorage* renderStorageSrc = mRenderStorageStack.last(); + WgRenderStorage* src = mRenderStorageStack.last(); mRenderStorageStack.pop(); - WgRenderStorage* renderStorageMsk = mRenderStorageStack.last(); + WgRenderStorage* msk = mRenderStorageStack.last(); mRenderStorageStack.pop(); - WgRenderStorage* renderStorageDst = mRenderStorageStack.last(); - - // get compose, blend and opacity settings - WgBindGroupCompositeMethod* composeMethod = mCompositeMethodPool.allocate(mContext, comp->method); - WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod); - WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, comp->opacity); - - // compose and blend - // dest = blend(dest, compose(src, msk, composeMethod), blendMethod, opacity) - mRenderStorageCopy.copy(mCommandEncoder, renderStorageDst); - renderStorageDst->compose(mContext, mCommandEncoder, - renderStorageSrc, renderStorageMsk, &mRenderStorageCopy, - composeMethod, blendMethod, opacity); - + WgRenderStorage* dst = mRenderStorageStack.last(); + // apply composition + mCompositor.compose(mCommandEncoder, src, msk, dst,comp->method); // back render targets to the pool - mRenderStoragePool.free(mContext, renderStorageSrc); - mRenderStoragePool.free(mContext, renderStorageMsk); - // begin last render pass - mRenderStorageStack.last()->beginRenderPass(mCommandEncoder, false); + mRenderStoragePool.free(mContext, src); + mRenderStoragePool.free(mContext, msk); } + // begin previous render pass + mCompositor.beginRenderPass(mCommandEncoder, mRenderStorageStack.last(), false); - // delete current compositor + // delete current compositor settings delete mCompositorStack.last(); mCompositorStack.pop(); diff --git a/src/renderer/wg_engine/tvgWgRenderer.h b/src/renderer/wg_engine/tvgWgRenderer.h old mode 100644 new mode 100755 index 799cb4e8..9d71805f --- a/src/renderer/wg_engine/tvgWgRenderer.h +++ b/src/renderer/wg_engine/tvgWgRenderer.h @@ -23,7 +23,7 @@ #ifndef _TVG_WG_RENDERER_H_ #define _TVG_WG_RENDERER_H_ -#include "tvgWgRenderTarget.h" +#include "tvgWgCompositor.h" class WgRenderer : public RenderMethod { @@ -61,37 +61,32 @@ private: ~WgRenderer(); void initialize(); void release(); - void clearDisposes(); - void renderClipPath(Array& clips); + void disposeObjects(); WGPUCommandEncoder mCommandEncoder{}; - WgRenderStorage mRenderStorageInterm; // intermidiate buffer to render - WgRenderStorage mRenderStorageCopy; // copy of destination target (blend and compostition) - WgRenderStorage mRenderStorageMask; // buffer to render mask - WgRenderStorage mRenderStorageRoot; // root render storage - WgRenderStorage mRenderStorageScreen; // storage with data after antializing - WgRenderStoragePool mRenderStoragePool; // pool to hold render tree storages - WgBindGroupOpacityPool mOpacityPool; // opacity, blend methods and composite methods pool - WgBindGroupBlendMethodPool mBlendMethodPool; - WgBindGroupCompositeMethodPool mCompositeMethodPool; - WgRenderDataShapePool mRenderDataShapePool; // render data shpes pool + WgRenderDataShapePool mRenderDataShapePool; // render tree stacks - Array mCompositorStack; + WgRenderStorage mStorageRoot; + Array mCompositorStack; Array mRenderStorageStack; + WgRenderStoragePool mRenderStoragePool; WgContext mContext; WgPipelines mPipelines; + WgCompositor mCompositor; + + // screen buffer + WGPUTexture mTexScreen{}; + WGPUTextureView mTexViewScreen{}; + WGPUBindGroup mBindGroupScreen{}; + Surface mTargetSurface; + BlendMethod mBlendMethod{}; RenderRegion mViewport{}; - BlendMethod mBlendMethod{}; // current blend method - - // disposed resources, they should be released on synced call. - struct { - Array renderDatas{}; - Key key; - } mDisposed; + Array mDisposeRenderDatas{}; + Key mDisposeKey{}; }; #endif /* _TVG_WG_RENDERER_H_ */ diff --git a/src/renderer/wg_engine/tvgWgShaderSrc.cpp b/src/renderer/wg_engine/tvgWgShaderSrc.cpp old mode 100644 new mode 100755 index 202c05d3..c9318ce1 --- a/src/renderer/wg_engine/tvgWgShaderSrc.cpp +++ b/src/renderer/wg_engine/tvgWgShaderSrc.cpp @@ -26,10 +26,10 @@ #define WG_SHADER_SOURCE(...) #__VA_ARGS__ //************************************************************************ -// shader pipeline fill +// graphics shader source: stencil //************************************************************************ -const char* cShaderSource_PipelineFill = WG_SHADER_SOURCE( +const char* cShaderSrc_Stencil = WG_SHADER_SOURCE( // vertex input struct VertexInput { @location(0) position: vec2f @@ -41,8 +41,8 @@ struct VertexOutput { }; // uniforms -@group(0) @binding(0) var uViewMat : mat4x4f; -@group(1) @binding(0) var uModelMat : mat4x4f; +@group(0) @binding(0) var uViewMat : mat4x4f; +@group(1) @binding(0) var uModelMat : mat4x4f; @vertex fn vs_main(in: VertexInput) -> VertexOutput { @@ -59,10 +59,10 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f { ); //************************************************************************ -// shader pipeline solid +// graphics shader source: solid //************************************************************************ -const char* cShaderSource_PipelineSolid = WG_SHADER_SOURCE( +const char* cShaderSrc_Solid = WG_SHADER_SOURCE( // vertex input struct VertexInput { @location(0) position: vec2f @@ -109,10 +109,10 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f { ); //************************************************************************ -// shader pipeline linear +// graphics shader source: linear //************************************************************************ -const char* cShaderSource_PipelineLinear = WG_SHADER_SOURCE( +const char* cShaderSrc_Linear = WG_SHADER_SOURCE( // vertex input struct VertexInput { @location(0) position: vec2f @@ -149,7 +149,7 @@ struct VertexOutput { // uniforms @group(0) @binding(0) var uViewMat : mat4x4f; @group(1) @binding(0) var uModelMat : mat4x4f; -@group(1) @binding(1) var uBlendSettings : BlendSettings; +@group(1) @binding(1) var uBlendSettings : BlendSettings; @group(2) @binding(0) var uLinearGradient : LinearGradient; @vertex @@ -212,10 +212,10 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f { ); //************************************************************************ -// shader pipeline radial +// graphics shader source: radial //************************************************************************ -const char* cShaderSource_PipelineRadial = WG_SHADER_SOURCE( +const char* cShaderSrc_Radial = WG_SHADER_SOURCE( // vertex input struct VertexInput { @location(0) position: vec2f @@ -309,10 +309,10 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f { ); //************************************************************************ -// cShaderSource_PipelineImage +// graphics shader source: image //************************************************************************ -const char* cShaderSource_PipelineImage = WG_SHADER_SOURCE( +const char* cShaderSrc_Image = WG_SHADER_SOURCE( // vertex input struct VertexInput { @location(0) position: vec2f, @@ -365,296 +365,329 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f { ); //************************************************************************ -// cShaderSource_PipelineComputeBlend +// compute shaders blend //************************************************************************ -const std::string strBlendShaderHeader = WG_SHADER_SOURCE( -@group(0) @binding(0) var imageSrc : texture_storage_2d; -@group(0) @binding(1) var imageDst : texture_storage_2d; -@group(0) @binding(2) var imageTrg : texture_storage_2d; -@group(1) @binding(0) var blendMethod : u32; -@group(2) @binding(0) var opacity : f32; -@compute @workgroup_size(8, 8) -fn cs_main( @builtin(global_invocation_id) id: vec3u) { - let texSize = textureDimensions(imageSrc); - if ((id.x >= texSize.x) || (id.y >= texSize.y)) { return; }; - - let colorSrc= textureLoad(imageSrc, id.xy); - if (colorSrc.a == 0.0) { return; } - let colorDst = textureLoad(imageDst, id.xy); - - var One: vec3f = vec3(1.0); - var So: f32 = opacity; - var Sc: vec3f = colorSrc.rgb; - var Sa: f32 = colorSrc.a; - var Dc: vec3f = colorDst.rgb; - var Da: f32 = colorDst.a; - var Rc: vec3f = colorDst.rgb; - var Ra: f32 = 1.0; -); - -const std::string strBlendMaskShaderHeader = WG_SHADER_SOURCE( -@group(0) @binding(0) var imageSrc : texture_storage_2d; -@group(0) @binding(1) var imageMsk : texture_storage_2d; -@group(0) @binding(2) var imageDst : texture_storage_2d; -@group(0) @binding(3) var imageTrg : texture_storage_2d; -@group(1) @binding(0) var blendMethod : u32; -@group(2) @binding(0) var opacity : f32; - -@compute @workgroup_size(8, 8) -fn cs_main( @builtin(global_invocation_id) id: vec3u) { - let texSize = textureDimensions(imageSrc); - if ((id.x >= texSize.x) || (id.y >= texSize.y)) { return; }; - - let colorMsk = textureLoad(imageMsk, id.xy); - if (colorMsk.a == 0.0) { return; } - let colorSrc = textureLoad(imageSrc, id.xy); - if (colorSrc.a == 0.0) { return; } - let colorDst = textureLoad(imageDst, id.xy); - - var One: vec3f = vec3(1.0); - var So: f32 = opacity; - var Sc: vec3f = colorSrc.rgb; - var Sa: f32 = colorSrc.a; - var Dc: vec3f = colorDst.rgb; - var Da: f32 = colorDst.a; - var Rc: vec3f = colorDst.rgb; - var Ra: f32 = 1.0; -); - -const std::string strBlendShaderPreConditionsGradient = WG_SHADER_SOURCE( - Sc = Sc + Dc.rgb * (1.0 - Sa); - Sa = Sa + Da * (1.0 - Sa); -); - -const std::string strBlendShaderPreConditionsImage = WG_SHADER_SOURCE( - Sc = Sc * So; - Sa = Sa * So; -); - -const std::string strBlendShaderBlendMethod = WG_SHADER_SOURCE( - switch blendMethod { - /* Normal */ case 0u: { - Rc = Sc + Dc * (1.0 - Sa); - Ra = Sa + Da * (1.0 - Sa); - } - /* Add */ case 1u: { Rc = Sc + Dc; } - /* Screen */ case 2u: { Rc = Sc + Dc - Sc * Dc; } - /* Multiply */ case 3u: { Rc = Sc * Dc; } - /* Overlay */ case 4u: { - Rc.r = select(1.0 - min(1.0, 2 * (1 - Sc.r) * (1 - Dc.r)), min(1.0, 2 * Sc.r * Dc.r), (Dc.r < 0.5)); - Rc.g = select(1.0 - min(1.0, 2 * (1 - Sc.g) * (1 - Dc.g)), min(1.0, 2 * Sc.g * Dc.g), (Dc.g < 0.5)); - Rc.b = select(1.0 - min(1.0, 2 * (1 - Sc.b) * (1 - Dc.b)), min(1.0, 2 * Sc.b * Dc.b), (Dc.b < 0.5)); - } - /* Difference */ case 5u: { Rc = abs(Dc - Sc); } - /* Exclusion */ case 6u: { Rc = min(One, Sc + Dc - min(One, 2 * Sc * Dc)); } - /* SrcOver */ case 7u: { Rc = Sc; Ra = Sa; } - /* Darken */ case 8u: { Rc = min(Sc, Dc); } - /* Lighten */ case 9u: { Rc = max(Sc, Dc); } - /* ColorDodge */ case 10u: { - Rc.r = select(Dc.r, (Dc.r * 255.0 / (255.0 - Sc.r * 255.0))/255.0, (1.0 - Sc.r > 0.0)); - Rc.g = select(Dc.g, (Dc.g * 255.0 / (255.0 - Sc.g * 255.0))/255.0, (1.0 - Sc.g > 0.0)); - Rc.b = select(Dc.b, (Dc.b * 255.0 / (255.0 - Sc.b * 255.0))/255.0, (1.0 - Sc.b > 0.0)); - } - /* ColorBurn */ case 11u: { - Rc.r = select(1.0 - Dc.r, (255.0 - (255.0 - Dc.r * 255.0) / (Sc.r * 255.0)) / 255.0, (Sc.r > 0.0)); - Rc.g = select(1.0 - Dc.g, (255.0 - (255.0 - Dc.g * 255.0) / (Sc.g * 255.0)) / 255.0, (Sc.g > 0.0)); - Rc.b = select(1.0 - Dc.b, (255.0 - (255.0 - Dc.b * 255.0) / (Sc.b * 255.0)) / 255.0, (Sc.b > 0.0)); - } - /* HardLight */ case 12u: { - Rc.r = select(1.0 - min(1.0, 2 * (1 - Sc.r) * (1 - Dc.r)), min(1.0, 2 * Sc.r * Dc.r), (Sc.r < 0.5)); - Rc.g = select(1.0 - min(1.0, 2 * (1 - Sc.g) * (1 - Dc.g)), min(1.0, 2 * Sc.g * Dc.g), (Sc.g < 0.5)); - Rc.b = select(1.0 - min(1.0, 2 * (1 - Sc.b) * (1 - Dc.b)), min(1.0, 2 * Sc.b * Dc.b), (Sc.b < 0.5)); - } - /* SoftLight */ case 13u: { Rc = min(One, (One - 2 * Sc) * Dc * Dc + 2.0 * Sc * Dc); } - default: { - Rc = Sc + Dc * (1.0 - Sa); - Ra = Sa + Da * (1.0 - Sa); - } - } -); - -const std::string strBlendShaderPostConditionsGradient = WG_SHADER_SOURCE( - // nothing -); - -const std::string strBlendShaderPostConditionsImage = WG_SHADER_SOURCE( - Rc = select(mix(Dc, Rc, Sa), Rc, blendMethod == 0); - Ra = select(mix(Da, Ra, Sa), Ra, blendMethod == 0); -); - -const std::string strBlendShaderFooter = WG_SHADER_SOURCE( - textureStore(imageTrg, id.xy, vec4(Rc, Ra)); -} -); - -// pipeline shader modules blend solid -const std::string strComputeBlendSolid = - strBlendShaderHeader + - strBlendShaderBlendMethod + - strBlendShaderFooter; -const char* cShaderSource_PipelineComputeBlendSolid = strComputeBlendSolid.c_str(); - -// pipeline shader modules blend gradient -const std::string strComputeBlendGradient = - strBlendShaderHeader + - strBlendShaderPreConditionsGradient + - strBlendShaderBlendMethod + - strBlendShaderPostConditionsGradient + - strBlendShaderFooter; -const char* cShaderSource_PipelineComputeBlendGradient = strComputeBlendGradient.c_str(); - -// pipeline shader modules blend image -const std::string strComputeBlendImage = - strBlendShaderHeader + - strBlendShaderPreConditionsImage + - strBlendShaderBlendMethod + - strBlendShaderPostConditionsImage + - strBlendShaderFooter; -const char* cShaderSource_PipelineComputeBlendImage = strComputeBlendImage.c_str(); - -// pipeline shader modules blend solid mask -const std::string strComputeBlendSolidMask = - strBlendMaskShaderHeader + - strBlendShaderBlendMethod + - strBlendShaderFooter; -const char* cShaderSource_PipelineComputeBlendSolidMask = strComputeBlendSolidMask.c_str(); - -// pipeline shader modules blend gradient mask -const std::string strComputeBlendGradientMask = - strBlendMaskShaderHeader + - strBlendShaderPreConditionsGradient + - strBlendShaderBlendMethod + - strBlendShaderPostConditionsGradient + - strBlendShaderFooter; -const char* cShaderSource_PipelineComputeBlendGradientMask = strComputeBlendGradientMask.c_str(); - -// pipeline shader modules blend image mask -const std::string strComputeBlendImageMask = - strBlendMaskShaderHeader + - strBlendShaderPreConditionsImage + - strBlendShaderBlendMethod + - strBlendShaderPostConditionsImage + - strBlendShaderFooter; -const char* cShaderSource_PipelineComputeBlendImageMask = strComputeBlendImageMask.c_str(); - -// pipeline shader modules clear -const char* cShaderSource_PipelineComputeClear = WG_SHADER_SOURCE( -@group(0) @binding(0) var imageDst : texture_storage_2d; - -@compute @workgroup_size(8, 8) -fn cs_main( @builtin(global_invocation_id) id: vec3u) { - textureStore(imageDst, id.xy, vec4f(0.0, 0.0, 0.0, 0.0)); -} -); - - -// pipeline shader modules copy -const char* cShaderSource_PipelineComputeCopy = WG_SHADER_SOURCE( -@group(0) @binding(0) var imageSrc : texture_storage_2d; -@group(1) @binding(0) var imageDst : texture_storage_2d; - -@compute @workgroup_size(8, 8) -fn cs_main( @builtin(global_invocation_id) id: vec3u) { - textureStore(imageDst, id.xy, textureLoad(imageSrc, id.xy)); -} -); - - -// pipeline shader modules compose -const char* cShaderSource_PipelineComputeMaskCompose = WG_SHADER_SOURCE( +const char* cShaderSrc_MergeMasks = WG_SHADER_SOURCE( @group(0) @binding(0) var imageMsk0 : texture_storage_2d; -@group(0) @binding(1) var imageMsk1 : texture_storage_2d; -@group(0) @binding(2) var imageTrg : texture_storage_2d; +@group(1) @binding(0) var imageMsk1 : texture_storage_2d; +@group(2) @binding(0) var imageTrg : texture_storage_2d; @compute @workgroup_size(8, 8) -fn cs_main( @builtin(global_invocation_id) id: vec3u) { +fn cs_main(@builtin(global_invocation_id) id: vec3u) { let colorMsk0 = textureLoad(imageMsk0, id.xy); let colorMsk1 = textureLoad(imageMsk1, id.xy); textureStore(imageTrg, id.xy, colorMsk0 * colorMsk1); } ); -// pipeline shader modules compose -const char* cShaderSource_PipelineComputeCompose = WG_SHADER_SOURCE( + +//************************************************************************ +// compute shaders blend +//************************************************************************ + +std::string cBlendHeader = WG_SHADER_SOURCE( @group(0) @binding(0) var imageSrc : texture_storage_2d; -@group(0) @binding(1) var imageMsk : texture_storage_2d; -@group(0) @binding(2) var imageDst : texture_storage_2d; -@group(0) @binding(3) var imageTrg : texture_storage_2d; -@group(1) @binding(0) var composeMethod : u32; -@group(2) @binding(0) var blendMethod : u32; -@group(3) @binding(0) var opacity : f32; +@group(1) @binding(0) var imageDst : texture_storage_2d; +@group(2) @binding(0) var imageTgt : texture_storage_2d; +@group(3) @binding(0) var So : 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; }; +fn cs_main(@builtin(global_invocation_id) id: vec3u) { + let colorSrc = textureLoad(imageSrc, id.xy); + if (colorSrc.a == 0.0) { return; } + let colorDst = textureLoad(imageDst, id.xy); + var Sc: vec3f = colorSrc.rgb; + var Sa: f32 = colorSrc.a; + let Dc: vec3f = colorDst.rgb; + let Da: f32 = colorDst.a; + var Rc: vec3f = colorDst.rgb; + var Ra: f32 = 1.0; +); + +const std::string cBlendPreProcess_Gradient = WG_SHADER_SOURCE( + Sc = Sc + Dc * (1.0 - Sa); + Sa = Sa + Da * (1.0 - Sa); +); +const std::string cBlendPreProcess_Image = WG_SHADER_SOURCE( + Sc = Sc * So; + Sa = Sa * So; +); + +const std::string cBlendEquation_Normal = WG_SHADER_SOURCE( + Rc = Sc + Dc * (1.0 - Sa); + Ra = Sa + Da * (1.0 - Sa); +); +const std::string cBlendEquation_Add = WG_SHADER_SOURCE( + Rc = Sc + Dc; +); +const std::string cBlendEquation_Screen = WG_SHADER_SOURCE( + Rc = Sc + Dc - Sc * Dc; +); +const std::string cBlendEquation_Multiply = WG_SHADER_SOURCE( + Rc = Sc * Dc; +); +const std::string cBlendEquation_Overlay = WG_SHADER_SOURCE( + Rc.r = select(1.0 - min(1.0, 2 * (1 - Sc.r) * (1 - Dc.r)), min(1.0, 2 * Sc.r * Dc.r), (Dc.r < 0.5)); + Rc.g = select(1.0 - min(1.0, 2 * (1 - Sc.g) * (1 - Dc.g)), min(1.0, 2 * Sc.g * Dc.g), (Dc.g < 0.5)); + Rc.b = select(1.0 - min(1.0, 2 * (1 - Sc.b) * (1 - Dc.b)), min(1.0, 2 * Sc.b * Dc.b), (Dc.b < 0.5)); +); +const std::string cBlendEquation_Difference = WG_SHADER_SOURCE( + Rc = abs(Dc - Sc); +); +const std::string cBlendEquation_Exclusion = WG_SHADER_SOURCE( + let One = vec3f(1.0, 1.0, 1.0); + Rc = min(One, Sc + Dc - min(One, 2 * Sc * Dc)); +); +const std::string cBlendEquation_SrcOver = WG_SHADER_SOURCE( + Rc = Sc; Ra = Sa; +); +const std::string cBlendEquation_Darken = WG_SHADER_SOURCE( + Rc = min(Sc, Dc); +); +const std::string cBlendEquation_Lighten = WG_SHADER_SOURCE( + Rc = max(Sc, Dc); +); +const std::string cBlendEquation_ColorDodge = WG_SHADER_SOURCE( + Rc.r = select(Dc.r, (Dc.r * 255.0 / (255.0 - Sc.r * 255.0))/255.0, (1.0 - Sc.r > 0.0)); + Rc.g = select(Dc.g, (Dc.g * 255.0 / (255.0 - Sc.g * 255.0))/255.0, (1.0 - Sc.g > 0.0)); + Rc.b = select(Dc.b, (Dc.b * 255.0 / (255.0 - Sc.b * 255.0))/255.0, (1.0 - Sc.b > 0.0)); +); +const std::string cBlendEquation_ColorBurn = WG_SHADER_SOURCE( + Rc.r = select(1.0 - Dc.r, (255.0 - (255.0 - Dc.r * 255.0) / (Sc.r * 255.0)) / 255.0, (Sc.r > 0.0)); + Rc.g = select(1.0 - Dc.g, (255.0 - (255.0 - Dc.g * 255.0) / (Sc.g * 255.0)) / 255.0, (Sc.g > 0.0)); + Rc.b = select(1.0 - Dc.b, (255.0 - (255.0 - Dc.b * 255.0) / (Sc.b * 255.0)) / 255.0, (Sc.b > 0.0)); +); +const std::string cBlendEquation_HardLight = WG_SHADER_SOURCE( + Rc.r = select(1.0 - min(1.0, 2 * (1 - Sc.r) * (1 - Dc.r)), min(1.0, 2 * Sc.r * Dc.r), (Sc.r < 0.5)); + Rc.g = select(1.0 - min(1.0, 2 * (1 - Sc.g) * (1 - Dc.g)), min(1.0, 2 * Sc.g * Dc.g), (Sc.g < 0.5)); + Rc.b = select(1.0 - min(1.0, 2 * (1 - Sc.b) * (1 - Dc.b)), min(1.0, 2 * Sc.b * Dc.b), (Sc.b < 0.5)); +); +const std::string cBlendEquation_SoftLight = WG_SHADER_SOURCE( + let One = vec3f(1.0, 1.0, 1.0); + Rc = min(One, (One - 2 * Sc) * Dc * Dc + 2.0 * Sc * Dc); +); + +const std::string cBlendPostProcess_Gradient = WG_SHADER_SOURCE( +); +const std::string cBlendPostProcess_Image = WG_SHADER_SOURCE( + Rc = mix(Dc, Rc, Sa); + Ra = mix(Da, Ra, Sa); +); + +const std::string cBlendFooter = WG_SHADER_SOURCE( + textureStore(imageTgt, id.xy, vec4f(Rc, Ra)); +} +); + +std::string blend_solid_Normal = cBlendHeader + cBlendEquation_Normal + cBlendFooter; +std::string blend_solid_Add = cBlendHeader + cBlendEquation_Add + cBlendFooter; +std::string blend_solid_Screen = cBlendHeader + cBlendEquation_Screen + cBlendFooter; +std::string blend_solid_Multiply = cBlendHeader + cBlendEquation_Multiply + cBlendFooter; +std::string blend_solid_Overlay = cBlendHeader + cBlendEquation_Overlay + cBlendFooter; +std::string blend_solid_Difference = cBlendHeader + cBlendEquation_Difference + cBlendFooter; +std::string blend_solid_Exclusion = cBlendHeader + cBlendEquation_Exclusion + cBlendFooter; +std::string blend_solid_SrcOver = cBlendHeader + cBlendEquation_SrcOver + cBlendFooter; +std::string blend_solid_Darken = cBlendHeader + cBlendEquation_Darken + cBlendFooter; +std::string blend_solid_Lighten = cBlendHeader + cBlendEquation_Lighten + cBlendFooter; +std::string blend_solid_ColorDodge = cBlendHeader + cBlendEquation_ColorDodge + cBlendFooter; +std::string blend_solid_ColorBurn = cBlendHeader + cBlendEquation_ColorBurn + cBlendFooter; +std::string blend_solid_HardLight = cBlendHeader + cBlendEquation_HardLight + cBlendFooter; +std::string blend_solid_SoftLight = cBlendHeader + cBlendEquation_SoftLight + cBlendFooter; + +const char* cShaderSrc_Blend_Solid[] { + blend_solid_Normal.c_str(), + blend_solid_Add.c_str(), + blend_solid_Screen.c_str(), + blend_solid_Multiply.c_str(), + blend_solid_Overlay.c_str(), + blend_solid_Difference.c_str(), + blend_solid_Exclusion.c_str(), + blend_solid_SrcOver.c_str(), + blend_solid_Darken.c_str(), + blend_solid_Lighten.c_str(), + blend_solid_ColorDodge.c_str(), + blend_solid_ColorBurn.c_str(), + blend_solid_HardLight.c_str(), + blend_solid_SoftLight.c_str() +}; + + +const std::string cBlendHeader_Gradient = cBlendHeader + cBlendPreProcess_Gradient; +const std::string cBlendFooter_Gradient = cBlendPostProcess_Gradient + cBlendFooter; +std::string blend_gradient_Normal = cBlendHeader_Gradient + cBlendEquation_Normal + cBlendFooter_Gradient; +std::string blend_gradient_Add = cBlendHeader_Gradient + cBlendEquation_Add + cBlendFooter_Gradient; +std::string blend_gradient_Screen = cBlendHeader_Gradient + cBlendEquation_Screen + cBlendFooter_Gradient; +std::string blend_gradient_Multiply = cBlendHeader_Gradient + cBlendEquation_Multiply + cBlendFooter_Gradient; +std::string blend_gradient_Overlay = cBlendHeader_Gradient + cBlendEquation_Overlay + cBlendFooter_Gradient; +std::string blend_gradient_Difference = cBlendHeader_Gradient + cBlendEquation_Difference + cBlendFooter_Gradient; +std::string blend_gradient_Exclusion = cBlendHeader_Gradient + cBlendEquation_Exclusion + cBlendFooter_Gradient; +std::string blend_gradient_SrcOver = cBlendHeader_Gradient + cBlendEquation_SrcOver + cBlendFooter_Gradient; +std::string blend_gradient_Darken = cBlendHeader_Gradient + cBlendEquation_Darken + cBlendFooter_Gradient; +std::string blend_gradient_Lighten = cBlendHeader_Gradient + cBlendEquation_Lighten + cBlendFooter_Gradient; +std::string blend_gradient_ColorDodge = cBlendHeader_Gradient + cBlendEquation_ColorDodge + cBlendFooter_Gradient; +std::string blend_gradient_ColorBurn = cBlendHeader_Gradient + cBlendEquation_ColorBurn + cBlendFooter_Gradient; +std::string blend_gradient_HardLight = cBlendHeader_Gradient + cBlendEquation_HardLight + cBlendFooter_Gradient; +std::string blend_gradient_SoftLight = cBlendHeader_Gradient + cBlendEquation_SoftLight + cBlendFooter_Gradient; + +const char* cShaderSrc_Blend_Gradient[] { + blend_gradient_Normal.c_str(), + blend_gradient_Add.c_str(), + blend_gradient_Screen.c_str(), + blend_gradient_Multiply.c_str(), + blend_gradient_Overlay.c_str(), + blend_gradient_Difference.c_str(), + blend_gradient_Exclusion.c_str(), + blend_gradient_SrcOver.c_str(), + blend_gradient_Darken.c_str(), + blend_gradient_Lighten.c_str(), + blend_gradient_ColorDodge.c_str(), + blend_gradient_ColorBurn.c_str(), + blend_gradient_HardLight.c_str(), + blend_gradient_SoftLight.c_str() +}; + +const std::string cBlendHeader_Image = cBlendHeader + cBlendPreProcess_Image; +const std::string cBlendFooter_Image = cBlendPostProcess_Image + cBlendFooter; +std::string blend_image_Normal = cBlendHeader_Image + cBlendEquation_Normal + cBlendFooter; +std::string blend_image_Add = cBlendHeader_Image + cBlendEquation_Add + cBlendFooter_Image; +std::string blend_image_Screen = cBlendHeader_Image + cBlendEquation_Screen + cBlendFooter_Image; +std::string blend_image_Multiply = cBlendHeader_Image + cBlendEquation_Multiply + cBlendFooter_Image; +std::string blend_image_Overlay = cBlendHeader_Image + cBlendEquation_Overlay + cBlendFooter_Image; +std::string blend_image_Difference = cBlendHeader_Image + cBlendEquation_Difference + cBlendFooter_Image; +std::string blend_image_Exclusion = cBlendHeader_Image + cBlendEquation_Exclusion + cBlendFooter_Image; +std::string blend_image_SrcOver = cBlendHeader_Image + cBlendEquation_SrcOver + cBlendFooter_Image; +std::string blend_image_Darken = cBlendHeader_Image + cBlendEquation_Darken + cBlendFooter_Image; +std::string blend_image_Lighten = cBlendHeader_Image + cBlendEquation_Lighten + cBlendFooter_Image; +std::string blend_image_ColorDodge = cBlendHeader_Image + cBlendEquation_ColorDodge + cBlendFooter_Image; +std::string blend_image_ColorBurn = cBlendHeader_Image + cBlendEquation_ColorBurn + cBlendFooter_Image; +std::string blend_image_HardLight = cBlendHeader_Image + cBlendEquation_HardLight + cBlendFooter_Image; +std::string blend_image_SoftLight = cBlendHeader_Image + cBlendEquation_SoftLight + cBlendFooter_Image; + +const char* cShaderSrc_Blend_Image[] { + blend_image_Normal.c_str(), + blend_image_Add.c_str(), + blend_image_Screen.c_str(), + blend_image_Multiply.c_str(), + blend_image_Overlay.c_str(), + blend_image_Difference.c_str(), + blend_image_Exclusion.c_str(), + blend_image_SrcOver.c_str(), + blend_image_Darken.c_str(), + blend_image_Lighten.c_str(), + blend_image_ColorDodge.c_str(), + blend_image_ColorBurn.c_str(), + blend_image_HardLight.c_str(), + blend_image_SoftLight.c_str() +}; + +//************************************************************************ +// compute shaders compose +//************************************************************************ + +std::string cComposeHeader = WG_SHADER_SOURCE( +@group(0) @binding(0) var imageSrc : texture_storage_2d; +@group(1) @binding(0) var imageMsk : texture_storage_2d; +@group(2) @binding(0) var imageDst : texture_storage_2d; +@group(3) @binding(0) var imageTgt : texture_storage_2d; + +@compute @workgroup_size(8, 8) +fn cs_main(@builtin(global_invocation_id) id: vec3u) { let colorSrc = textureLoad(imageSrc, id.xy); let colorMsk = textureLoad(imageMsk, id.xy); let colorDst = textureLoad(imageDst, id.xy); - var One: vec3f = vec3(1.0); - let luma: f32 = dot(colorMsk.xyz, vec3f(0.2125, 0.7154, 0.0721)); - var So : f32 = opacity; - var Mc : vec3f = colorMsk.rgb; - var Ma : f32 = colorMsk.a; - var MSc : vec3f = colorSrc.rgb; - var MSa : f32 = colorSrc.a; - var Sc : vec3f = colorSrc.rgb; - var Sa : f32 = colorSrc.a; - var Dc : vec3f = colorDst.rgb; - var Da : f32 = colorDst.a; - var Rc : vec3f = colorDst.rgb; - var Ra : f32 = 1.0; + var Sc : vec3f = colorSrc.rgb; + var Sa : f32 = colorSrc.a; + var Mc : vec3f = colorMsk.rgb; + var Ma : f32 = colorMsk.a; + var Dc : vec3f = colorDst.rgb; + var Da : f32 = colorDst.a; + var Rc : vec3f = colorDst.rgb; + var Ra : f32 = colorDst.a; +); - if (composeMethod <= /* InvLumaMask */5u) { - switch composeMethod { - /* AlphaMask */ case 2u: { Sc = MSc * Ma; Sa = MSa * Ma; } - /* InvAlphaMask */ case 3u: { Sc = MSc * (1.0 - Ma); Sa = MSa * (1.0 - Ma); } - /* LumaMask */ case 4u: { Sc = MSc * luma; Sa = MSa * luma; } - /* InvLumaMask */ case 5u: { Sc = MSc * (1.0-luma); Sa = MSa * (1.0-luma); } - default: { Sc = MSc; Sa = MSa; } - } - Rc = Sc + Dc * (1.0 - Sa); - Ra = Sa + Da * (1.0 - Sa); - } else { - Sc = Dc; - switch composeMethod { - /* AddMask */ case 6u: { Sa = MSa + Ma * (1.0 - MSa); } - /* SubtractMask */ case 7u: { Sa = MSa * (1.0 - Ma); } - /* IntersectMask */ case 8u: { Sa = MSa * Ma; } - /* DifferenceMask */ case 9u: { Sa = MSa * (1.0 - Ma) + Ma * (1.0 - MSa); } - default: { Rc = Sc; Ra = Sa; } - } - Rc = Sc; - Ra = Sa; - } +std::string cComposeEquation_None = WG_SHADER_SOURCE( + Rc = Dc; + Ra = Da; +); +std::string cComposeEquation_ClipPath = WG_SHADER_SOURCE( + Rc = Sc * Ma + Dc * (1.0 - Sa * Ma); + Ra = Sa * Ma + Da * (1.0 - Sa * Ma); +); - textureStore(imageTrg, id.xy, vec4f(Rc, Ra)); +std::string cComposeEquation_AlphaMask = WG_SHADER_SOURCE( + Rc = Sc * Ma + Dc * (1.0 - Sa * Ma); + Ra = Sa * Ma + Da * (1.0 - Sa * Ma); +); +std::string cComposeEquation_InvAlphaMask = WG_SHADER_SOURCE( + Rc = Sc * (1.0 - Ma) + Dc * (1.0 - Sa * (1.0 - Ma)); + Ra = Sa * (1.0 - Ma) + Da * (1.0 - Sa * (1.0 - Ma)); +); +std::string cComposeEquation_LumaMask = WG_SHADER_SOURCE( + let luma: f32 = dot(Mc, vec3f(0.2125, 0.7154, 0.0721)); + Rc = Sc * luma + Dc * (1.0 - Sa * luma); + Ra = Sa * luma + Da * (1.0 - Sa * luma); +); +std::string cComposeEquation_InvLumaMask = WG_SHADER_SOURCE( + let luma: f32 = dot(Mc, vec3f(0.2125, 0.7154, 0.0721)); + Rc = Sc * (1.0 - luma) + Dc * (1.0 - Sa * (1.0 - luma)); + Ra = Sa * (1.0 - luma) + Da * (1.0 - Sa * (1.0 - luma)); +); +std::string cComposeEquation_AddMask = WG_SHADER_SOURCE( + Rc = Sc; + Ra = Sa + Ma * (1.0 - Sa); +); +std::string cComposeEquation_SubtractMask = WG_SHADER_SOURCE( + Rc = Sc; + Ra = Sa * (1.0 - Ma); +); +std::string cComposeEquation_IntersectMask = WG_SHADER_SOURCE( + Rc = Sc; + Ra = Sa * Ma; +); +std::string cComposeEquation_DifferenceMask = WG_SHADER_SOURCE( + Rc = Sc; + Ra = Sa * (1.0 - Ma) + Ma * (1.0 - Sa); +); + +std::string cComposeFooter = WG_SHADER_SOURCE( + textureStore(imageTgt, id.xy, vec4f(Rc, Ra)); } ); -// pipeline shader modules anti-aliasing -const char* cShaderSource_PipelineComputeAntiAlias = WG_SHADER_SOURCE( +std::string compose_None = cComposeHeader + cComposeEquation_None + cComposeFooter; +std::string compose_ClipPath = cComposeHeader + cComposeEquation_ClipPath + cComposeFooter; +std::string compose_AlphaMask = cComposeHeader + cComposeEquation_AlphaMask + cComposeFooter; +std::string compose_InvAlphaMask = cComposeHeader + cComposeEquation_InvAlphaMask + cComposeFooter; +std::string compose_LumaMask = cComposeHeader + cComposeEquation_LumaMask + cComposeFooter; +std::string compose_InvLumaMask = cComposeHeader + cComposeEquation_InvLumaMask + cComposeFooter; +std::string compose_AddMask = cComposeHeader + cComposeEquation_AddMask + cComposeFooter; +std::string compose_SubtractMask = cComposeHeader + cComposeEquation_SubtractMask + cComposeFooter; +std::string compose_IntersectMask = cComposeHeader + cComposeEquation_IntersectMask + cComposeFooter; +std::string compose_DifferenceMask = cComposeHeader + cComposeEquation_DifferenceMask + cComposeFooter; + +const char* cShaderSrc_Compose[10] { + compose_None.c_str(), + compose_ClipPath.c_str(), + compose_AlphaMask.c_str(), + compose_InvAlphaMask.c_str(), + compose_LumaMask.c_str(), + compose_InvLumaMask.c_str(), + compose_AddMask.c_str(), + compose_SubtractMask.c_str(), + compose_IntersectMask.c_str(), + compose_DifferenceMask.c_str(), +}; + +const char* cShaderSrc_Copy = WG_SHADER_SOURCE( @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 texSizeSrc = textureDimensions(imageSrc); - let texSizeDst = textureDimensions(imageDst); - if ((id.x >= texSizeDst.x) || (id.y >= texSizeDst.y)) { return; }; - - let samples = u32(texSizeSrc.x / texSizeDst.x); - var color = vec4f(0); - for (var i = 0u; i < samples; i++) { - for (var j = 0u; j < samples; j++) { - color += textureLoad(imageSrc, vec2(id.x * samples + j, id.y * samples + i)); - } - } - - textureStore(imageDst, id.xy, color / f32(samples * samples)); +fn cs_main(@builtin(global_invocation_id) id: vec3u) { + textureStore(imageDst, id.xy, textureLoad(imageSrc, id.xy)); } -); +); \ No newline at end of file diff --git a/src/renderer/wg_engine/tvgWgShaderSrc.h b/src/renderer/wg_engine/tvgWgShaderSrc.h old mode 100644 new mode 100755 index 0a051a2b..51706cd6 --- a/src/renderer/wg_engine/tvgWgShaderSrc.h +++ b/src/renderer/wg_engine/tvgWgShaderSrc.h @@ -20,37 +20,23 @@ * SOFTWARE. */ -#pragma once - #ifndef _TVG_WG_SHADER_SRC_H_ #define _TVG_WG_SHADER_SRC_H_ -//***************************************************************************** -// render shader modules -//***************************************************************************** +// graphics shader sources +extern const char* cShaderSrc_Stencil; +extern const char* cShaderSrc_Solid; +extern const char* cShaderSrc_Linear; +extern const char* cShaderSrc_Radial; +extern const char* cShaderSrc_Image; -// pipeline shader modules fill -extern const char* cShaderSource_PipelineFill; -extern const char* cShaderSource_PipelineSolid; -extern const char* cShaderSource_PipelineLinear; -extern const char* cShaderSource_PipelineRadial; -extern const char* cShaderSource_PipelineImage; +// compute shader sources: blend, compose and merge path +extern const char* cShaderSrc_MergeMasks; +extern const char* cShaderSrc_Blend_Solid[14]; +extern const char* cShaderSrc_Blend_Solid[14]; +extern const char* cShaderSrc_Blend_Gradient[14]; +extern const char* cShaderSrc_Blend_Image[14]; +extern const char* cShaderSrc_Compose[10]; +extern const char* cShaderSrc_Copy; -//***************************************************************************** -// compute shader modules -//***************************************************************************** - -// pipeline shader modules clear, compose and blend -extern const char* cShaderSource_PipelineComputeClear; -extern const char* cShaderSource_PipelineComputeCopy; -extern const char* cShaderSource_PipelineComputeBlendSolid; -extern const char* cShaderSource_PipelineComputeBlendGradient; -extern const char* cShaderSource_PipelineComputeBlendImage; -extern const char* cShaderSource_PipelineComputeBlendSolidMask; -extern const char* cShaderSource_PipelineComputeBlendGradientMask; -extern const char* cShaderSource_PipelineComputeBlendImageMask; -extern const char* cShaderSource_PipelineComputeMaskCompose; -extern const char* cShaderSource_PipelineComputeCompose; -extern const char* cShaderSource_PipelineComputeAntiAlias; - -#endif // _TVG_WG_SHADER_SRC_H_ +#endif // _TVG_WG_SHEDER_SRC_H_ diff --git a/src/renderer/wg_engine/tvgWgShaderTypes.cpp b/src/renderer/wg_engine/tvgWgShaderTypes.cpp old mode 100644 new mode 100755 diff --git a/src/renderer/wg_engine/tvgWgShaderTypes.h b/src/renderer/wg_engine/tvgWgShaderTypes.h old mode 100644 new mode 100755 index 3c033533..f9cc1015 --- a/src/renderer/wg_engine/tvgWgShaderTypes.h +++ b/src/renderer/wg_engine/tvgWgShaderTypes.h @@ -23,7 +23,7 @@ #ifndef _TVG_WG_SHADER_TYPES_H_ #define _TVG_WG_SHADER_TYPES_H_ -#include "tvgWgCommon.h" +#include "tvgRender.h" /////////////////////////////////////////////////////////////////////////////// // shader types