diff --git a/src/renderer/wg_engine/tvgWgBindGroups.cpp b/src/renderer/wg_engine/tvgWgBindGroups.cpp index 07e76697..b051a6ab 100644 --- a/src/renderer/wg_engine/tvgWgBindGroups.cpp +++ b/src/renderer/wg_engine/tvgWgBindGroups.cpp @@ -33,11 +33,15 @@ WGPUBindGroupLayout WgBindGroupRadialGradient::layout = nullptr; WGPUBindGroupLayout WgBindGroupPicture::layout = nullptr; // composition and blending properties groups WGPUBindGroupLayout WgBindGroupTexture::layout = nullptr; -WGPUBindGroupLayout WgBindGroupTextureStorageRgba::layout = nullptr; -WGPUBindGroupLayout WgBindGroupTextureStorageBgra::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 WgBindGroupTexComposeBlend::layout = nullptr; +WGPUBindGroupLayout WgBindGroupTexCompose::layout = nullptr; WGPUBindGroupLayout WgBindGroupOpacity::layout = nullptr; WGPUBindGroupLayout WgBindGroupBlendMethod::layout = nullptr; WGPUBindGroupLayout WgBindGroupCompositeMethod::layout = nullptr; @@ -338,14 +342,14 @@ void WgBindGroupTexture::release() } /////////////////////////////////////////////////////////////////////////////// -// WgBindGroupTextureStorageRgba +// WgBindGroupTextureStorageRgbaRO /////////////////////////////////////////////////////////////////////////////// -WGPUBindGroupLayout WgBindGroupTextureStorageRgba::getLayout(WGPUDevice device) +WGPUBindGroupLayout WgBindGroupTextureStorageRgbaRO::getLayout(WGPUDevice device) { if (layout) return layout; const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { - makeBindGroupLayoutEntryStorage(0, WGPUStorageTextureAccess_ReadWrite, WGPUTextureFormat_RGBA8Unorm) + makeBindGroupLayoutEntryStorage(0, WGPUStorageTextureAccess_ReadOnly, WGPUTextureFormat_RGBA8Unorm) }; layout = createBindGroupLayout(device, bindGroupLayoutEntries, 1); assert(layout); @@ -353,13 +357,13 @@ WGPUBindGroupLayout WgBindGroupTextureStorageRgba::getLayout(WGPUDevice device) } -void WgBindGroupTextureStorageRgba::releaseLayout() +void WgBindGroupTextureStorageRgbaRO::releaseLayout() { releaseBindGroupLayout(layout); } -void WgBindGroupTextureStorageRgba::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexture) +void WgBindGroupTextureStorageRgbaRO::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexture) { release(); const WGPUBindGroupEntry bindGroupEntries[] { @@ -370,16 +374,91 @@ void WgBindGroupTextureStorageRgba::initialize(WGPUDevice device, WGPUQueue queu } -void WgBindGroupTextureStorageRgba::release() +void WgBindGroupTextureStorageRgbaRO::release() { releaseBindGroup(mBindGroup); } /////////////////////////////////////////////////////////////////////////////// -// WgBindGroupTextureStorageBgra +// WgBindGroupTextureStorageRgbaRO /////////////////////////////////////////////////////////////////////////////// -WGPUBindGroupLayout WgBindGroupTextureStorageBgra::getLayout(WGPUDevice device) +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[] { @@ -391,13 +470,13 @@ WGPUBindGroupLayout WgBindGroupTextureStorageBgra::getLayout(WGPUDevice device) } -void WgBindGroupTextureStorageBgra::releaseLayout() +void WgBindGroupTextureStorageBgraWO::releaseLayout() { releaseBindGroupLayout(layout); } -void WgBindGroupTextureStorageBgra::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexture) +void WgBindGroupTextureStorageBgraWO::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexture) { release(); const WGPUBindGroupEntry bindGroupEntries[] { @@ -408,7 +487,7 @@ void WgBindGroupTextureStorageBgra::initialize(WGPUDevice device, WGPUQueue queu } -void WgBindGroupTextureStorageBgra::release() +void WgBindGroupTextureStorageBgraWO::release() { releaseBindGroup(mBindGroup); } @@ -454,16 +533,16 @@ void WgBindGroupTextureSampled::release() } /////////////////////////////////////////////////////////////////////////////// -// WgBindGroupTexComposeBlend +// WgBindGroupTexBlend /////////////////////////////////////////////////////////////////////////////// -WGPUBindGroupLayout WgBindGroupTexComposeBlend::getLayout(WGPUDevice device) +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_ReadWrite, WGPUTextureFormat_RGBA8Unorm) + makeBindGroupLayoutEntryStorage(2, WGPUStorageTextureAccess_WriteOnly, WGPUTextureFormat_RGBA8Unorm) }; layout = createBindGroupLayout(device, bindGroupLayoutEntries, 3); assert(layout); @@ -471,13 +550,13 @@ WGPUBindGroupLayout WgBindGroupTexComposeBlend::getLayout(WGPUDevice device) } -void WgBindGroupTexComposeBlend::releaseLayout() +void WgBindGroupTexBlend::releaseLayout() { releaseBindGroupLayout(layout); } -void WgBindGroupTexComposeBlend::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexSrc, WGPUTextureView uTexMsk, WGPUTextureView uTexDst) +void WgBindGroupTexBlend::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexSrc, WGPUTextureView uTexMsk, WGPUTextureView uTexDst) { release(); const WGPUBindGroupEntry bindGroupEntries[] { @@ -490,7 +569,51 @@ void WgBindGroupTexComposeBlend::initialize(WGPUDevice device, WGPUQueue queue, } -void WgBindGroupTexComposeBlend::release() +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); } @@ -504,9 +627,10 @@ WGPUBindGroupLayout WgBindGroupTexMaskCompose::getLayout(WGPUDevice device) if (layout) return layout; const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { makeBindGroupLayoutEntryStorage(0, WGPUStorageTextureAccess_ReadOnly, WGPUTextureFormat_RGBA8Unorm), - makeBindGroupLayoutEntryStorage(1, WGPUStorageTextureAccess_ReadWrite, WGPUTextureFormat_RGBA8Unorm) + makeBindGroupLayoutEntryStorage(1, WGPUStorageTextureAccess_ReadOnly, WGPUTextureFormat_RGBA8Unorm), + makeBindGroupLayoutEntryStorage(2, WGPUStorageTextureAccess_WriteOnly, WGPUTextureFormat_RGBA8Unorm) }; - layout = createBindGroupLayout(device, bindGroupLayoutEntries, 2); + layout = createBindGroupLayout(device, bindGroupLayoutEntries, 3); assert(layout); return layout; } @@ -518,14 +642,15 @@ void WgBindGroupTexMaskCompose::releaseLayout() } -void WgBindGroupTexMaskCompose::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexMsk0, WGPUTextureView uTexMsk1) +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, 2); + mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 3); assert(mBindGroup); } @@ -535,6 +660,50 @@ 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 /////////////////////////////////////////////////////////////////////////////// diff --git a/src/renderer/wg_engine/tvgWgBindGroups.h b/src/renderer/wg_engine/tvgWgBindGroups.h index 0d9b33e1..d1837b91 100644 --- a/src/renderer/wg_engine/tvgWgBindGroups.h +++ b/src/renderer/wg_engine/tvgWgBindGroups.h @@ -119,7 +119,7 @@ struct WgBindGroupTexture : public WgBindGroup }; // @group(0 or 1) -struct WgBindGroupTextureStorageRgba : public WgBindGroup +struct WgBindGroupTextureStorageRgbaWO : public WgBindGroup { static WGPUBindGroupLayout layout; static WGPUBindGroupLayout getLayout(WGPUDevice device); @@ -131,7 +131,7 @@ struct WgBindGroupTextureStorageRgba : public WgBindGroup }; // @group(0 or 1) -struct WgBindGroupTextureStorageBgra : public WgBindGroup +struct WgBindGroupTextureStorageRgbaRO : public WgBindGroup { static WGPUBindGroupLayout layout; static WGPUBindGroupLayout getLayout(WGPUDevice device); @@ -142,6 +142,37 @@ struct WgBindGroupTextureStorageBgra : public WgBindGroup 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 { @@ -156,7 +187,22 @@ struct WgBindGroupTextureSampled : public WgBindGroup }; // @group(0) -struct WgBindGroupTexComposeBlend : public WgBindGroup +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); @@ -165,7 +211,8 @@ struct WgBindGroupTexComposeBlend : public WgBindGroup void initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexSrc, WGPUTextureView uTexMsk, - WGPUTextureView uTexDst); + WGPUTextureView uTexDst, + WGPUTextureView uTexTrg); void release(); }; @@ -179,7 +226,24 @@ struct WgBindGroupTexMaskCompose : public WgBindGroup void initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexMsk0, - WGPUTextureView uTexMsk1); + 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(); }; diff --git a/src/renderer/wg_engine/tvgWgCommon.cpp b/src/renderer/wg_engine/tvgWgCommon.cpp index bb0ebfc2..d84cfa45 100644 --- a/src/renderer/wg_engine/tvgWgCommon.cpp +++ b/src/renderer/wg_engine/tvgWgCommon.cpp @@ -21,6 +21,9 @@ */ #include "tvgWgCommon.h" +#ifdef __EMSCRIPTEN__ +#include +#endif #include //***************************************************************************** @@ -50,12 +53,13 @@ void WgContext::initialize(WGPUInstance instance, WGPUSurface surface) }; // request adapter wgpuInstanceRequestAdapter(instance, &requestAdapterOptions, onAdapterRequestEnded, &adapter); + #ifdef __EMSCRIPTEN__ + while (!adapter) emscripten_sleep(10); + #endif assert(adapter); - // adapter enumerate features + // get feature names size_t featuresCount = wgpuAdapterEnumerateFeatures(adapter, featureNames); - wgpuAdapterGetProperties(adapter, &adapterProperties); - wgpuAdapterGetLimits(adapter, &supportedLimits); // request device WGPUDeviceDescriptor deviceDesc{}; @@ -76,6 +80,9 @@ void WgContext::initialize(WGPUInstance instance, WGPUSurface surface) }; // request device wgpuAdapterRequestDevice(adapter, &deviceDesc, onDeviceRequestEnded, &device); + #ifdef __EMSCRIPTEN__ + while (!device) emscripten_sleep(10); + #endif assert(device); // on device error function @@ -535,8 +542,10 @@ WGPUShaderModule WgPipeline::createShaderModule(WGPUDevice device, const char* c WGPUShaderModuleDescriptor shaderModuleDesc{}; shaderModuleDesc.nextInChain = &shaderModuleWGSLDesc.chain; shaderModuleDesc.label = label; + #ifndef __EMSCRIPTEN__ shaderModuleDesc.hintCount = 0; shaderModuleDesc.hints = nullptr; + #endif return wgpuDeviceCreateShaderModule(device, &shaderModuleDesc); } diff --git a/src/renderer/wg_engine/tvgWgPipelines.cpp b/src/renderer/wg_engine/tvgWgPipelines.cpp index 4e0ce4de..1d87fec4 100644 --- a/src/renderer/wg_engine/tvgWgPipelines.cpp +++ b/src/renderer/wg_engine/tvgWgPipelines.cpp @@ -298,17 +298,18 @@ void WgPipelineImage::initialize(WGPUDevice device, WgPipelineBlendType blendTyp // compute pipelines //************************************************************************ -void WgPipelineClear::initialize(WGPUDevice device) +void WgPipelineCopy::initialize(WGPUDevice device) { // bind groups and layouts WGPUBindGroupLayout bindGroupLayouts[] = { - WgBindGroupTextureStorageRgba::getLayout(device) + WgBindGroupTextureStorageRgbaRO::getLayout(device), + WgBindGroupTextureStorageRgbaWO::getLayout(device) }; // shader source and labels - auto shaderSource = cShaderSource_PipelineComputeClear; - auto shaderLabel = "The compute shader clear"; - auto pipelineLabel = "The compute pipeline clear"; + auto shaderSource = cShaderSource_PipelineComputeCopy; + auto shaderLabel = "The compute shader copy"; + auto pipelineLabel = "The compute pipeline copy"; // allocate all pipeline handles allocate(device, @@ -321,8 +322,7 @@ void WgPipelineBlend::initialize(WGPUDevice device, const char *shaderSource) { // bind groups and layouts WGPUBindGroupLayout bindGroupLayouts[] = { - WgBindGroupTextureStorageRgba::getLayout(device), - WgBindGroupTextureStorageRgba::getLayout(device), + WgBindGroupTexBlend::getLayout(device), WgBindGroupBlendMethod::getLayout(device), WgBindGroupOpacity::getLayout(device) }; @@ -342,7 +342,7 @@ void WgPipelineBlendMask::initialize(WGPUDevice device, const char *shaderSource { // bind groups and layouts WGPUBindGroupLayout bindGroupLayouts[] = { - WgBindGroupTexComposeBlend::getLayout(device), + WgBindGroupTexBlendMask::getLayout(device), WgBindGroupBlendMethod::getLayout(device), WgBindGroupOpacity::getLayout(device) }; @@ -381,7 +381,7 @@ void WgPipelineCompose::initialize(WGPUDevice device) { // bind groups and layouts WGPUBindGroupLayout bindGroupLayouts[] = { - WgBindGroupTexComposeBlend::getLayout(device), + WgBindGroupTexCompose::getLayout(device), WgBindGroupCompositeMethod::getLayout(device), WgBindGroupBlendMethod::getLayout(device), WgBindGroupOpacity::getLayout(device) @@ -403,8 +403,8 @@ void WgPipelineAntiAliasing::initialize(WGPUDevice device) { // bind groups and layouts WGPUBindGroupLayout bindGroupLayouts[] = { - WgBindGroupTextureStorageRgba::getLayout(device), - WgBindGroupTextureStorageBgra::getLayout(device) + WgBindGroupTextureStorageRgbaRO::getLayout(device), + WgBindGroupTextureStorageBgraWO::getLayout(device) }; // shader source and labels @@ -436,7 +436,7 @@ void WgPipelines::initialize(WgContext& context) image[type].initialize(context.device, (WgPipelineBlendType)type); } // compute pipelines - computeClear.initialize(context.device); + computeCopy.initialize(context.device); computeBlendSolid.initialize(context.device, cShaderSource_PipelineComputeBlendSolid); computeBlendGradient.initialize(context.device, cShaderSource_PipelineComputeBlendGradient); computeBlendImage.initialize(context.device, cShaderSource_PipelineComputeBlendImage); @@ -453,11 +453,14 @@ void WgPipelines::initialize(WgContext& context) void WgPipelines::release() { + WgBindGroupTexCompose::releaseLayout(); WgBindGroupTexMaskCompose::releaseLayout(); - WgBindGroupTexComposeBlend::releaseLayout(); - WgBindGroupTextureSampled::releaseLayout(); - WgBindGroupTextureStorageBgra::releaseLayout(); - WgBindGroupTextureStorageRgba::releaseLayout(); + WgBindGroupTexBlendMask::releaseLayout(); + WgBindGroupTexBlend::releaseLayout(); + WgBindGroupTextureStorageBgraRO::releaseLayout(); + WgBindGroupTextureStorageBgraWO::releaseLayout(); + WgBindGroupTextureStorageRgbaRO::releaseLayout(); + WgBindGroupTextureStorageRgbaWO::releaseLayout(); WgBindGroupTexture::releaseLayout(); WgBindGroupOpacity::releaseLayout(); WgBindGroupPicture::releaseLayout(); @@ -469,14 +472,13 @@ void WgPipelines::release() // compute pipelines computeAntiAliasing.release(); computeCompose.release(); - computeMaskCompose.release(); computeBlendImageMask.release(); computeBlendGradientMask.release(); computeBlendSolidMask.release(); computeBlendImage.release(); computeBlendGradient.release(); computeBlendSolid.release(); - computeClear.release(); + computeCopy.release(); // fill pipelines for (uint8_t type = (uint8_t)WgPipelineBlendType::SrcOver; type <= (uint8_t)WgPipelineBlendType::Custom; type++) { image[type].release(); diff --git a/src/renderer/wg_engine/tvgWgPipelines.h b/src/renderer/wg_engine/tvgWgPipelines.h index 2a94a9cd..092cd6f5 100644 --- a/src/renderer/wg_engine/tvgWgPipelines.h +++ b/src/renderer/wg_engine/tvgWgPipelines.h @@ -129,13 +129,14 @@ struct WgPipelineImage: public WgRenderPipeline // compute pipelines //***************************************************************************** -struct WgPipelineClear: public WgComputePipeline +struct WgPipelineCopy: public WgComputePipeline { void initialize(WGPUDevice device) override; - void use(WGPUComputePassEncoder encoder, WgBindGroupTextureStorageRgba& groupTexDst) + void use(WGPUComputePassEncoder encoder, WgBindGroupTextureStorageRgbaRO& groupTexSrc, WgBindGroupTextureStorageRgbaWO& groupTexDst) { set(encoder); - groupTexDst.set(encoder, 0); + groupTexSrc.set(encoder, 0); + groupTexDst.set(encoder, 1); } }; @@ -144,13 +145,12 @@ struct WgPipelineBlend: public WgComputePipeline { void initialize(WGPUDevice device) override { assert(false); }; void initialize(WGPUDevice device, const char *shaderSource); - void use(WGPUComputePassEncoder encoder, WgBindGroupTextureStorageRgba& groupTexSrc, WgBindGroupTextureStorageRgba& groupTexDst, WgBindGroupBlendMethod& blendMethod, WgBindGroupOpacity& groupOpacity) + void use(WGPUComputePassEncoder encoder, WgBindGroupTexBlend& groupTexBlend, WgBindGroupBlendMethod& blendMethod, WgBindGroupOpacity& groupOpacity) { set(encoder); - groupTexSrc.set(encoder, 0); - groupTexDst.set(encoder, 1); - blendMethod.set(encoder, 2); - groupOpacity.set(encoder, 3); + groupTexBlend.set(encoder, 0); + blendMethod.set(encoder, 1); + groupOpacity.set(encoder, 2); } }; @@ -159,10 +159,10 @@ struct WgPipelineBlendMask: public WgComputePipeline { void initialize(WGPUDevice device) override { assert(false); }; void initialize(WGPUDevice device, const char *shaderSource); - void use(WGPUComputePassEncoder encoder, WgBindGroupTexComposeBlend& groupTexs, WgBindGroupBlendMethod& blendMethod, WgBindGroupOpacity& groupOpacity) + void use(WGPUComputePassEncoder encoder, WgBindGroupTexBlendMask& groupTexBlendMask, WgBindGroupBlendMethod& blendMethod, WgBindGroupOpacity& groupOpacity) { set(encoder); - groupTexs.set(encoder, 0); + groupTexBlendMask.set(encoder, 0); blendMethod.set(encoder, 1); groupOpacity.set(encoder, 2); } @@ -172,10 +172,10 @@ struct WgPipelineBlendMask: public WgComputePipeline struct WgPipelineMaskCompose: public WgComputePipeline { void initialize(WGPUDevice device) override; - void use(WGPUComputePassEncoder encoder, WgBindGroupTexMaskCompose& groupTexs) + void use(WGPUComputePassEncoder encoder, WgBindGroupTexMaskCompose& groupTexMaskCompose) { set(encoder); - groupTexs.set(encoder, 0); + groupTexMaskCompose.set(encoder, 0); } }; @@ -183,10 +183,10 @@ struct WgPipelineMaskCompose: public WgComputePipeline struct WgPipelineCompose: public WgComputePipeline { void initialize(WGPUDevice device) override; - void use(WGPUComputePassEncoder encoder, WgBindGroupTexComposeBlend& groupTexs, WgBindGroupCompositeMethod& groupComposeMethod, WgBindGroupBlendMethod& groupBlendMethod, WgBindGroupOpacity& groupOpacity) + void use(WGPUComputePassEncoder encoder, WgBindGroupTexCompose& groupTexCompose, WgBindGroupCompositeMethod& groupComposeMethod, WgBindGroupBlendMethod& groupBlendMethod, WgBindGroupOpacity& groupOpacity) { set(encoder); - groupTexs.set(encoder, 0); + groupTexCompose.set(encoder, 0); groupComposeMethod.set(encoder, 1); groupBlendMethod.set(encoder, 2); groupOpacity.set(encoder, 3); @@ -197,7 +197,7 @@ struct WgPipelineCompose: public WgComputePipeline struct WgPipelineAntiAliasing: public WgComputePipeline { void initialize(WGPUDevice device) override; - void use(WGPUComputePassEncoder encoder, WgBindGroupTextureStorageRgba& groupTexSrc, WgBindGroupTextureStorageBgra& groupTexDst) + void use(WGPUComputePassEncoder encoder, WgBindGroupTextureStorageRgbaRO& groupTexSrc, WgBindGroupTextureStorageBgraWO& groupTexDst) { set(encoder); groupTexSrc.set(encoder, 0); @@ -222,7 +222,7 @@ struct WgPipelines WgPipelineRadial radial[3]; WgPipelineImage image[3]; // compute pipelines - WgPipelineClear computeClear; + WgPipelineCopy computeCopy; WgPipelineBlend computeBlendSolid; WgPipelineBlend computeBlendGradient; WgPipelineBlend computeBlendImage; diff --git a/src/renderer/wg_engine/tvgWgRenderTarget.cpp b/src/renderer/wg_engine/tvgWgRenderTarget.cpp index 21314580..09cc5ea4 100644 --- a/src/renderer/wg_engine/tvgWgRenderTarget.cpp +++ b/src/renderer/wg_engine/tvgWgRenderTarget.cpp @@ -54,10 +54,12 @@ assert(texViewColor); assert(texViewStencil); // initialize bind group for blitting - if (format == WGPUTextureFormat_RGBA8Unorm) - bindGroupTexStorageRgba.initialize(context.device, context.queue, texViewColor); + if (format == WGPUTextureFormat_RGBA8Unorm) { + bindGroupTexStorageRgbaRO.initialize(context.device, context.queue, texViewColor); + bindGroupTexStorageRgbaWO.initialize(context.device, context.queue, texViewColor); + } if (format == WGPUTextureFormat_BGRA8Unorm) - bindGroupTexStorageBgra.initialize(context.device, context.queue, texViewColor); + bindGroupTexStorageBgraWO.initialize(context.device, context.queue, texViewColor); // initialize window binding groups WgShaderTypeMat4x4f viewMat(w, h); mBindGroupCanvas.initialize(context.device, context.queue, viewMat); @@ -69,7 +71,9 @@ void WgRenderStorage::release(WgContext& context) { mRenderPassEncoder = nullptr; mBindGroupCanvas.release(); - bindGroupTexStorageRgba.release(); + bindGroupTexStorageBgraWO.release(); + bindGroupTexStorageRgbaWO.release(); + bindGroupTexStorageRgbaRO.release(); context.releaseTextureView(texViewStencil); context.releaseTextureView(texViewColor); context.releaseTexture(texStencil); @@ -210,29 +214,24 @@ void WgRenderStorage::drawPictureClipPath(WgContext& context, WgRenderDataPictur } -void WgRenderStorage::clear(WGPUCommandEncoder commandEncoder) -{ - assert(commandEncoder); - WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder); - mPipelines->computeClear.use(computePassEncoder, bindGroupTexStorageRgba); - dispatchWorkgroups(computePassEncoder); - endComputePass(computePassEncoder); -} - void WgRenderStorage::blend( WgContext& context, WGPUCommandEncoder commandEncoder, WgPipelineBlend* pipeline, - WgRenderStorage* targetSrc, + WgRenderStorage* texSrc, + WgRenderStorage* texDst, WgBindGroupBlendMethod* blendMethod, WgBindGroupOpacity* opacity) { assert(commandEncoder); - assert(targetSrc); + 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, targetSrc->bindGroupTexStorageRgba, bindGroupTexStorageRgba, *blendMethod, *opacity); + pipeline->use(computePassEncoder, texBlend, *blendMethod, *opacity); dispatchWorkgroups(computePassEncoder); endComputePass(computePassEncoder); } @@ -244,31 +243,33 @@ void WgRenderStorage::blendMask( WgPipelineBlendMask* pipeline, WgRenderStorage* texMsk, WgRenderStorage* texSrc, + WgRenderStorage* texDst, WgBindGroupBlendMethod* blendMethod, WgBindGroupOpacity* opacity) { assert(commandEncoder); assert(texSrc); assert(texMsk); - WgBindGroupTexComposeBlend composeBlend; - composeBlend.initialize(context.device, context.queue, texSrc->texViewColor, texMsk->texViewColor, texViewColor); + WgBindGroupTexBlendMask texBlendMask; + texBlendMask.initialize(context.device, context.queue, texSrc->texViewColor, texMsk->texViewColor, texDst->texViewColor, texViewColor); WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder); - pipeline->use(computePassEncoder, composeBlend, *blendMethod, *opacity); + pipeline->use(computePassEncoder, texBlendMask, *blendMethod, *opacity); dispatchWorkgroups(computePassEncoder); endComputePass(computePassEncoder); - composeBlend.release(); + texBlendMask.release(); }; void WgRenderStorage::maskCompose( WgContext& context, WGPUCommandEncoder commandEncoder, - WgRenderStorage* texMsk0) + WgRenderStorage* texMsk0, + WgRenderStorage* texMsk1) { assert(commandEncoder); assert(texMsk0); WgBindGroupTexMaskCompose maskCompose; - maskCompose.initialize(context.device, context.queue, texMsk0->texViewColor, texViewColor); + maskCompose.initialize(context.device, context.queue, texMsk0->texViewColor, texMsk1->texViewColor, texViewColor); WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder); mPipelines->computeMaskCompose.use(computePassEncoder, maskCompose); dispatchWorkgroups(computePassEncoder); @@ -282,6 +283,7 @@ void WgRenderStorage::compose( WGPUCommandEncoder commandEncoder, WgRenderStorage* texSrc, WgRenderStorage* texMsk, + WgRenderStorage* texDst, WgBindGroupCompositeMethod* composeMethod, WgBindGroupBlendMethod* blendMethod, WgBindGroupOpacity* opacity) @@ -289,13 +291,13 @@ void WgRenderStorage::compose( assert(commandEncoder); assert(texSrc); assert(texMsk); - WgBindGroupTexComposeBlend composeBlend; - composeBlend.initialize(context.device, context.queue, texSrc->texViewColor, texMsk->texViewColor, texViewColor); + WgBindGroupTexCompose texCompose; + texCompose.initialize(context.device, context.queue, texSrc->texViewColor, texMsk->texViewColor, texDst->texViewColor, texViewColor); WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder); - mPipelines->computeCompose.use(computePassEncoder, composeBlend, *composeMethod, *blendMethod, *opacity); + mPipelines->computeCompose.use(computePassEncoder, texCompose, *composeMethod, *blendMethod, *opacity); dispatchWorkgroups(computePassEncoder); endComputePass(computePassEncoder); - composeBlend.release(); + texCompose.release(); } @@ -304,11 +306,20 @@ void WgRenderStorage::antialias(WGPUCommandEncoder commandEncoder, WgRenderStora assert(commandEncoder); assert(targetSrc); WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder); - mPipelines->computeAntiAliasing.use(computePassEncoder, targetSrc->bindGroupTexStorageRgba, bindGroupTexStorageBgra); + 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) { @@ -323,8 +334,8 @@ void WgRenderStorage::beginRenderPass(WGPUCommandEncoder commandEncoder, bool cl // render pass depth stencil attachment WGPURenderPassDepthStencilAttachment depthStencilAttachment{}; depthStencilAttachment.view = texViewStencil; - depthStencilAttachment.depthLoadOp = WGPULoadOp_Load; - depthStencilAttachment.depthStoreOp = WGPUStoreOp_Discard; + depthStencilAttachment.depthLoadOp = WGPULoadOp_Undefined; + depthStencilAttachment.depthStoreOp = WGPUStoreOp_Undefined; depthStencilAttachment.depthClearValue = 1.0f; depthStencilAttachment.depthReadOnly = false; depthStencilAttachment.stencilLoadOp = WGPULoadOp_Clear; @@ -334,6 +345,9 @@ void WgRenderStorage::beginRenderPass(WGPUCommandEncoder commandEncoder, bool cl // 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 }; diff --git a/src/renderer/wg_engine/tvgWgRenderTarget.h b/src/renderer/wg_engine/tvgWgRenderTarget.h index 7d9ba590..70e4b87f 100644 --- a/src/renderer/wg_engine/tvgWgRenderTarget.h +++ b/src/renderer/wg_engine/tvgWgRenderTarget.h @@ -36,8 +36,9 @@ public: WGPUTexture texStencil{}; WGPUTextureView texViewColor{}; WGPUTextureView texViewStencil{}; - WgBindGroupTextureStorageRgba bindGroupTexStorageRgba; - WgBindGroupTextureStorageBgra bindGroupTexStorageBgra; + WgBindGroupTextureStorageRgbaRO bindGroupTexStorageRgbaRO; + WgBindGroupTextureStorageRgbaWO bindGroupTexStorageRgbaWO; + WgBindGroupTextureStorageBgraWO bindGroupTexStorageBgraWO; uint32_t samples{}; uint32_t width{}; uint32_t height{}; @@ -54,12 +55,12 @@ public: void renderPicture(WgContext& context, WgRenderDataPicture* renderData, WgPipelineBlendType blendType); void renderClipPath(WgContext& context, WgRenderDataPaint* renderData); - void clear(WGPUCommandEncoder commandEncoder); void blend( WgContext& context, WGPUCommandEncoder commandEncoder, WgPipelineBlend* pipeline, WgRenderStorage* targetSrc, + WgRenderStorage* targetDst, WgBindGroupBlendMethod* blendMethod, WgBindGroupOpacity* opacity); void blendMask( @@ -68,21 +69,25 @@ public: WgPipelineBlendMask* pipeline, WgRenderStorage* texMsk, WgRenderStorage* texSrc, + WgRenderStorage* texDst, WgBindGroupBlendMethod* blendMethod, WgBindGroupOpacity* opacity); void maskCompose( WgContext& context, WGPUCommandEncoder commandEncoder, - WgRenderStorage* texMsk0); + 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); diff --git a/src/renderer/wg_engine/tvgWgRenderer.cpp b/src/renderer/wg_engine/tvgWgRenderer.cpp index 4640a506..40961688 100644 --- a/src/renderer/wg_engine/tvgWgRenderer.cpp +++ b/src/renderer/wg_engine/tvgWgRenderer.cpp @@ -68,6 +68,7 @@ void WgRenderer::release() mOpacityPool.release(mContext); mRenderStorageRoot.release(mContext); mRenderStorageMask.release(mContext); + mRenderStorageCopy.release(mContext); mRenderStorageScreen.release(mContext); mRenderStorageInterm.release(mContext); mPipelines.release(); @@ -190,7 +191,8 @@ void WgRenderer::renderClipPath(Array& clips) mRenderStorageInterm.beginRenderPass(mCommandEncoder, true); mRenderStorageInterm.renderClipPath(mContext, clips[i]); mRenderStorageInterm.endRenderPass(); - mRenderStorageMask.maskCompose(mContext, mCommandEncoder, &mRenderStorageInterm); + mRenderStorageCopy.copy(mCommandEncoder, &mRenderStorageMask); + mRenderStorageMask.maskCompose(mContext, mCommandEncoder, &mRenderStorageInterm, &mRenderStorageCopy); } } @@ -222,8 +224,9 @@ bool WgRenderer::renderShape(RenderData data) 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, blendMethod, opacity); + pipeline, &mRenderStorageMask, &mRenderStorageInterm, &mRenderStorageCopy, blendMethod, opacity); // restore current render pass renderStorage->beginRenderPass(mCommandEncoder, false); // use hardware blend @@ -243,8 +246,9 @@ bool WgRenderer::renderShape(RenderData data) 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, blendMethod, opacity); + pipeline, &mRenderStorageInterm, &mRenderStorageCopy, blendMethod, opacity); // restore current render pass renderStorage->beginRenderPass(mCommandEncoder, false); } @@ -277,8 +281,9 @@ bool WgRenderer::renderImage(RenderData data) 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, blendMethod, opacity); + pipeline, &mRenderStorageMask, &mRenderStorageInterm, &mRenderStorageCopy, blendMethod, opacity); // restore current render pass renderStorage->beginRenderPass(mCommandEncoder, false); // use hardware blend @@ -296,8 +301,9 @@ bool WgRenderer::renderImage(RenderData data) 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, blendMethod, opacity); + pipeline, &mRenderStorageInterm, &mRenderStorageCopy, blendMethod, opacity); // restore current render pass renderStorage->beginRenderPass(mCommandEncoder, false); } @@ -418,12 +424,18 @@ bool WgRenderer::target(WGPUInstance instance, WGPUSurface surface, uint32_t w, 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); 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); return true; @@ -469,9 +481,10 @@ bool WgRenderer::endComposite(TVG_UNUSED Compositor* cmp) // 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, blendMethod, opacity); + renderStorageSrc, &mRenderStorageCopy, blendMethod, opacity); // back render targets to the pool mRenderStoragePool.free(mContext, renderStorageSrc); @@ -495,8 +508,9 @@ bool WgRenderer::endComposite(TVG_UNUSED Compositor* cmp) // compose and blend // dest = blend(dest, compose(src, msk, composeMethod), blendMethod, opacity) + mRenderStorageCopy.copy(mCommandEncoder, renderStorageDst); renderStorageDst->compose(mContext, mCommandEncoder, - renderStorageSrc, renderStorageMsk, + renderStorageSrc, renderStorageMsk, &mRenderStorageCopy, composeMethod, blendMethod, opacity); // back render targets to the pool diff --git a/src/renderer/wg_engine/tvgWgRenderer.h b/src/renderer/wg_engine/tvgWgRenderer.h index 193872d4..799cb4e8 100644 --- a/src/renderer/wg_engine/tvgWgRenderer.h +++ b/src/renderer/wg_engine/tvgWgRenderer.h @@ -66,6 +66,7 @@ private: 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 diff --git a/src/renderer/wg_engine/tvgWgShaderSrc.cpp b/src/renderer/wg_engine/tvgWgShaderSrc.cpp index 51767c20..202c05d3 100644 --- a/src/renderer/wg_engine/tvgWgShaderSrc.cpp +++ b/src/renderer/wg_engine/tvgWgShaderSrc.cpp @@ -369,10 +369,11 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f { //************************************************************************ const std::string strBlendShaderHeader = WG_SHADER_SOURCE( -@group(0) @binding(0) var imageSrc : texture_storage_2d; -@group(1) @binding(0) var imageDst : texture_storage_2d; -@group(2) @binding(0) var blendMethod : u32; -@group(3) @binding(0) var opacity : f32; +@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) { @@ -396,7 +397,8 @@ fn cs_main( @builtin(global_invocation_id) id: vec3u) { 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(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; @@ -483,7 +485,7 @@ const std::string strBlendShaderPostConditionsImage = WG_SHADER_SOURCE( ); const std::string strBlendShaderFooter = WG_SHADER_SOURCE( - textureStore(imageDst, id.xy, vec4(Rc, Ra)); + textureStore(imageTrg, id.xy, vec4(Rc, Ra)); } ); @@ -539,7 +541,7 @@ const char* cShaderSource_PipelineComputeBlendImageMask = strComputeBlendImageMa // pipeline shader modules clear const char* cShaderSource_PipelineComputeClear = WG_SHADER_SOURCE( -@group(0) @binding(0) var imageDst : texture_storage_2d; +@group(0) @binding(0) var imageDst : texture_storage_2d; @compute @workgroup_size(8, 8) fn cs_main( @builtin(global_invocation_id) id: vec3u) { @@ -547,16 +549,30 @@ fn cs_main( @builtin(global_invocation_id) id: vec3u) { } ); + +// 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( @group(0) @binding(0) var imageMsk0 : texture_storage_2d; -@group(0) @binding(1) var imageMsk1 : texture_storage_2d; +@group(0) @binding(1) var imageMsk1 : texture_storage_2d; +@group(0) @binding(2) var imageTrg : texture_storage_2d; @compute @workgroup_size(8, 8) fn cs_main( @builtin(global_invocation_id) id: vec3u) { let colorMsk0 = textureLoad(imageMsk0, id.xy); let colorMsk1 = textureLoad(imageMsk1, id.xy); - textureStore(imageMsk1, id.xy, colorMsk0 * colorMsk1); + textureStore(imageTrg, id.xy, colorMsk0 * colorMsk1); } ); @@ -564,7 +580,8 @@ fn cs_main( @builtin(global_invocation_id) id: vec3u) { const char* cShaderSource_PipelineComputeCompose = 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(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; @@ -615,13 +632,13 @@ fn cs_main( @builtin(global_invocation_id) id: vec3u) { Ra = Sa; } - textureStore(imageDst, id.xy, vec4f(Rc, Ra)); + textureStore(imageTrg, id.xy, vec4f(Rc, Ra)); } ); // pipeline shader modules anti-aliasing const char* cShaderSource_PipelineComputeAntiAlias = WG_SHADER_SOURCE( -@group(0) @binding(0) var imageSrc : texture_storage_2d; +@group(0) @binding(0) var imageSrc : texture_storage_2d; @group(1) @binding(0) var imageDst : texture_storage_2d; @compute @workgroup_size(8, 8) diff --git a/src/renderer/wg_engine/tvgWgShaderSrc.h b/src/renderer/wg_engine/tvgWgShaderSrc.h index eb5f97bb..0a051a2b 100644 --- a/src/renderer/wg_engine/tvgWgShaderSrc.h +++ b/src/renderer/wg_engine/tvgWgShaderSrc.h @@ -42,6 +42,7 @@ extern const char* cShaderSource_PipelineImage; // 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;