wg_engine: introduced compute pipeline entities

introduces posibility to create compute pipelines
does not affect functionality

src/renderer/wg_engine/tvgWgBindGroups.cpp

@@ -22,15 +22,20 @@
 
 #include "tvgWgBindGroups.h"
 
-
+// canvas information group
 WGPUBindGroupLayout WgBindGroupCanvas::layout = nullptr;
+// paint object information group
 WGPUBindGroupLayout WgBindGroupPaint::layout = nullptr;
+// fill properties information groups
 WGPUBindGroupLayout WgBindGroupSolidColor::layout = nullptr;
 WGPUBindGroupLayout WgBindGroupLinearGradient::layout = nullptr;
 WGPUBindGroupLayout WgBindGroupRadialGradient::layout = nullptr;
 WGPUBindGroupLayout WgBindGroupPicture::layout = nullptr;
+// composition and blending properties gropus
 WGPUBindGroupLayout WgBindGroupOpacity::layout = nullptr;
-WGPUBindGroupLayout WgBindGroupBlit::layout = nullptr;
+WGPUBindGroupLayout WgBindGroupTexture::layout = nullptr;
+WGPUBindGroupLayout WgBindGroupStorageTexture::layout = nullptr;
+WGPUBindGroupLayout WgBindGroupTextureSampled::layout = nullptr;
 
 
 WGPUBindGroupLayout WgBindGroupCanvas::getLayout(WGPUDevice device)
@@ -227,7 +232,7 @@ WGPUBindGroupLayout WgBindGroupPicture::getLayout(WGPUDevice device)
     if (layout) return layout;
     const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] {
         makeBindGroupLayoutEntrySampler(0),
-        makeBindGroupLayoutEntryTextureView(1)
+        makeBindGroupLayoutEntryTexture(1)
     };
     layout = createBindGroupLayout(device, bindGroupLayoutEntries, 2);
     assert(layout);
@@ -303,12 +308,82 @@ void WgBindGroupOpacity::release()
 }
 
 
-WGPUBindGroupLayout WgBindGroupBlit::getLayout(WGPUDevice device)
+WGPUBindGroupLayout WgBindGroupTexture::getLayout(WGPUDevice device)
+{
+    if (layout) return layout;
+    const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] {
+        makeBindGroupLayoutEntryTexture(0)
+    };
+    layout = createBindGroupLayout(device, bindGroupLayoutEntries, 1);
+    assert(layout);
+    return layout;
+}
+
+
+void WgBindGroupTexture::releaseLayout()
+{
+    releaseBindGroupLayout(layout);
+}
+
+
+void WgBindGroupTexture::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexture)
+{
+    release();
+    const WGPUBindGroupEntry bindGroupEntries[] {
+        makeBindGroupEntryTextureView(0, uTexture)
+    };
+    mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 1);
+    assert(mBindGroup);
+}
+
+
+void WgBindGroupTexture::release()
+{
+    releaseBindGroup(mBindGroup);
+}
+
+
+WGPUBindGroupLayout WgBindGroupStorageTexture::getLayout(WGPUDevice device)
+{
+    if (layout) return layout;
+    const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] {
+        makeBindGroupLayoutEntryStorageTexture(0)
+    };
+    layout = createBindGroupLayout(device, bindGroupLayoutEntries, 1);
+    assert(layout);
+    return layout;
+}
+
+
+void WgBindGroupStorageTexture::releaseLayout()
+{
+    releaseBindGroupLayout(layout);
+}
+
+
+void WgBindGroupStorageTexture::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexture)
+{
+    release();
+    const WGPUBindGroupEntry bindGroupEntries[] {
+        makeBindGroupEntryTextureView(0, uTexture)
+    };
+    mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 1);
+    assert(mBindGroup);
+}
+
+
+void WgBindGroupStorageTexture::release()
+{
+    releaseBindGroup(mBindGroup);
+}
+
+
+WGPUBindGroupLayout WgBindGroupTextureSampled::getLayout(WGPUDevice device)
 {
     if (layout) return layout;
     const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] {
         makeBindGroupLayoutEntrySampler(0),
-        makeBindGroupLayoutEntryTextureView(1)
+        makeBindGroupLayoutEntryTexture(1)
     };
     layout = createBindGroupLayout(device, bindGroupLayoutEntries, 2);
     assert(layout);
@@ -316,13 +391,13 @@ WGPUBindGroupLayout WgBindGroupBlit::getLayout(WGPUDevice device)
 }
 
 
-void WgBindGroupBlit::releaseLayout()
+void WgBindGroupTextureSampled::releaseLayout()
 {
     releaseBindGroupLayout(layout);
 }
 
 
-void WgBindGroupBlit::initialize(WGPUDevice device, WGPUQueue queue, WGPUSampler uSampler, WGPUTextureView uTexture)
+void WgBindGroupTextureSampled::initialize(WGPUDevice device, WGPUQueue queue, WGPUSampler uSampler, WGPUTextureView uTexture)
 {
     release();
     const WGPUBindGroupEntry bindGroupEntries[] {
@@ -334,49 +409,40 @@ void WgBindGroupBlit::initialize(WGPUDevice device, WGPUQueue queue, WGPUSampler
 }
 
 
-void WgBindGroupBlit::release()
+void WgBindGroupTextureSampled::release()
 {
     releaseBindGroup(mBindGroup);
 }
 
 //************************************************************************
-// bind groups pools
+// bind group pools
 //************************************************************************
 
-WgBindGroupOpacity* WgBindGroupOpacityPool::allocate(WgContext& context, uint32_t opacity)
+void WgBindGroupOpacityPool::initialize(WgContext& context)
 {
-    WgBindGroupOpacity* bindGroup{};
+    memset(mPool, 0x00, sizeof(mPool));
-    if (mPool.count == 0) {
+}
-        bindGroup = new WgBindGroupOpacity;
+
-        bindGroup->initialize(context.device, context.queue, opacity);
+
-        mUsed.push(bindGroup);
+void WgBindGroupOpacityPool::release(WgContext& context)
-        mList.push(bindGroup);
+{
-    } else {
+    for (uint32_t i = 0; i < 256; i++) {
-        bindGroup = mPool.last();
+        if (mPool[i]) {
-        bindGroup->update(context.device, context.queue, opacity);
+            mPool[i]->release();
-        mUsed.push(bindGroup);
+            delete mPool[i];
-        mPool.pop();
+            mPool[i] = nullptr;
+        }
     }
-    return bindGroup;
 }
 
 
-void WgBindGroupOpacityPool::reset()
+WgBindGroupOpacity* WgBindGroupOpacityPool::allocate(WgContext& context, uint8_t opacity)
 {
-    for (uint32_t i = 0; i < mUsed.count; i++)
+    WgBindGroupOpacity* bindGroupOpacity = mPool[opacity];
-        mPool.push(mUsed[i]);
+    if (!bindGroupOpacity) {
-    mUsed.clear();
+        bindGroupOpacity = new WgBindGroupOpacity;
-}
+        bindGroupOpacity->initialize(context.device, context.queue, opacity);
-
+        mPool[opacity] = bindGroupOpacity;
-
-void WgBindGroupOpacityPool::release()
-{
-    for (uint32_t i = 0; i < mList.count; i++) {
-        mList[i]->release();
-        delete mList[i];
     }
-    mList.clear();
+    return bindGroupOpacity;
-    mUsed.clear();
-    mPool.clear();
 }
-

src/renderer/wg_engine/tvgWgBindGroups.h

@@ -39,7 +39,6 @@ struct WgBindGroupCanvas : public WgBindGroup
     void release();
 };
 
-
 // @group(1)
 struct WgBindGroupPaint : public WgBindGroup
 {
@@ -56,7 +55,7 @@ struct WgBindGroupPaint : public WgBindGroup
 };
 
 // @group(2)
-struct WgBindGroupSolidColor : public  WgBindGroup
+struct WgBindGroupSolidColor : public WgBindGroup
 {
     static WGPUBindGroupLayout layout;
     static WGPUBindGroupLayout getLayout(WGPUDevice device);
@@ -69,7 +68,7 @@ struct WgBindGroupSolidColor : public  WgBindGroup
 };
 
 // @group(2)
-struct WgBindGroupLinearGradient : public  WgBindGroup
+struct WgBindGroupLinearGradient : public WgBindGroup
 {
     static WGPUBindGroupLayout layout;
     static WGPUBindGroupLayout getLayout(WGPUDevice device);
@@ -82,7 +81,7 @@ struct WgBindGroupLinearGradient : public  WgBindGroup
 };
 
 // @group(2)
-struct WgBindGroupRadialGradient : public  WgBindGroup
+struct WgBindGroupRadialGradient : public WgBindGroup
 {
     static WGPUBindGroupLayout layout;
     static WGPUBindGroupLayout getLayout(WGPUDevice device);
@@ -95,7 +94,7 @@ struct WgBindGroupRadialGradient : public  WgBindGroup
 };
 
 // @group(2)
-struct WgBindGroupPicture : public  WgBindGroup
+struct WgBindGroupPicture : public WgBindGroup
 {
     static WGPUBindGroupLayout layout;
     static WGPUBindGroupLayout getLayout(WGPUDevice device);
@@ -108,7 +107,7 @@ struct WgBindGroupPicture : public  WgBindGroup
 };
 
 // @group(1 or 2)
-struct WgBindGroupOpacity : public  WgBindGroup
+struct WgBindGroupOpacity : public WgBindGroup
 {
     static WGPUBindGroupLayout layout;
     static WGPUBindGroupLayout getLayout(WGPUDevice device);
@@ -121,7 +120,31 @@ struct WgBindGroupOpacity : public  WgBindGroup
 };
 
 // @group(0 or 1)
-struct WgBindGroupBlit : public  WgBindGroup
+struct WgBindGroupTexture : public WgBindGroup
+{
+    static WGPUBindGroupLayout layout;
+    static WGPUBindGroupLayout getLayout(WGPUDevice device);
+    static void releaseLayout();
+
+    void initialize(WGPUDevice device, WGPUQueue queue,
+                    WGPUTextureView uTexture);
+    void release();
+};
+
+// @group(0 or 1)
+struct WgBindGroupStorageTexture : public WgBindGroup
+{
+    static WGPUBindGroupLayout layout;
+    static WGPUBindGroupLayout getLayout(WGPUDevice device);
+    static void releaseLayout();
+
+    void initialize(WGPUDevice device, WGPUQueue queue,
+                    WGPUTextureView uTexture);
+    void release();
+};
+
+// @group(0 or 1)
+struct WgBindGroupTextureSampled : public WgBindGroup
 {
     static WGPUBindGroupLayout layout;
     static WGPUBindGroupLayout getLayout(WGPUDevice device);
@@ -140,13 +163,11 @@ struct WgBindGroupBlit : public  WgBindGroup
 class WgBindGroupOpacityPool
 {
 private:
-    Array<WgBindGroupOpacity*> mList;
+    WgBindGroupOpacity* mPool[256];
-    Array<WgBindGroupOpacity*> mPool;
-    Array<WgBindGroupOpacity*> mUsed;
 public:
-    WgBindGroupOpacity* allocate(WgContext& context, uint32_t opacity);
+    void initialize(WgContext& context);
-    void reset();
+    void release(WgContext& context);
-    void release();
+    WgBindGroupOpacity* allocate(WgContext& context, uint8_t opacity);
 };
 
 #endif // _TVG_WG_BIND_GROUPS_H_

src/renderer/wg_engine/tvgWgCommon.cpp

@@ -60,7 +60,7 @@ void WgContext::initialize()
     WGPUDeviceDescriptor deviceDesc{};
     deviceDesc.nextInChain = nullptr;
     deviceDesc.label = "The device";
-    deviceDesc.requiredFeaturesCount = featuresCount;
+    deviceDesc.requiredFeatureCount = featuresCount;
     deviceDesc.requiredFeatures = featureNames;
     deviceDesc.requiredLimits = nullptr;
     deviceDesc.defaultQueue.nextInChain = nullptr;
@@ -115,6 +115,54 @@ void WgContext::executeCommandEncoder(WGPUCommandEncoder commandEncoder)
 }
 
 
+WGPUTexture WgContext::createTexture2d(WGPUTextureUsage usage, WGPUTextureFormat format, uint32_t width, uint32_t height, char const * label) {
+    WGPUTextureDescriptor textureDesc{};
+    textureDesc.nextInChain = nullptr;
+    textureDesc.label = label;
+    textureDesc.usage = usage;
+    textureDesc.dimension = WGPUTextureDimension_2D;
+    textureDesc.size = { width, height, 1 };
+    textureDesc.format = format;
+    textureDesc.mipLevelCount = 1;
+    textureDesc.sampleCount = 1;
+    textureDesc.viewFormatCount = 0;
+    textureDesc.viewFormats = nullptr;
+    return wgpuDeviceCreateTexture(device, &textureDesc);
+}
+
+
+WGPUTextureView WgContext::createTextureView2d(WGPUTexture texture, WGPU_NULLABLE char const * label)
+{
+    WGPUTextureViewDescriptor textureViewDescColor{};
+    textureViewDescColor.nextInChain = nullptr;
+    textureViewDescColor.label = label;
+    textureViewDescColor.format = wgpuTextureGetFormat(texture);
+    textureViewDescColor.dimension = WGPUTextureViewDimension_2D;
+    textureViewDescColor.baseMipLevel = 0;
+    textureViewDescColor.mipLevelCount = 1;
+    textureViewDescColor.baseArrayLayer = 0;
+    textureViewDescColor.arrayLayerCount = 1;
+    textureViewDescColor.aspect = WGPUTextureAspect_All;
+    return wgpuTextureCreateView(texture, &textureViewDescColor);
+};
+
+
+void WgContext::releaseTexture(WGPUTexture& texture) {
+    if (texture) {
+        wgpuTextureDestroy(texture);
+        wgpuTextureRelease(texture);
+        texture = nullptr;
+    }
+    
+}
+
+
+void WgContext::releaseTextureView(WGPUTextureView& textureView) {
+    if (textureView) wgpuTextureViewRelease(textureView);
+    textureView = nullptr;
+}
+
+
 //*****************************************************************************
 // bind group
 //*****************************************************************************
@@ -125,6 +173,12 @@ void WgBindGroup::set(WGPURenderPassEncoder encoder, uint32_t groupIndex)
 }
 
 
+void WgBindGroup::set(WGPUComputePassEncoder encoder, uint32_t groupIndex)
+{
+    wgpuComputePassEncoderSetBindGroup(encoder, groupIndex, mBindGroup, 0, nullptr);
+}
+
+
 WGPUBindGroupEntry WgBindGroup::makeBindGroupEntryBuffer(uint32_t binding, WGPUBuffer buffer)
 {
     WGPUBindGroupEntry bindGroupEntry{};
@@ -193,12 +247,12 @@ WGPUBindGroupLayoutEntry WgBindGroup::makeBindGroupLayoutEntrySampler(uint32_t b
 }
 
 
-WGPUBindGroupLayoutEntry WgBindGroup::makeBindGroupLayoutEntryTextureView(uint32_t binding)
+WGPUBindGroupLayoutEntry WgBindGroup::makeBindGroupLayoutEntryTexture(uint32_t binding)
 {
     WGPUBindGroupLayoutEntry bindGroupLayoutEntry{};
     bindGroupLayoutEntry.nextInChain = nullptr;
     bindGroupLayoutEntry.binding = binding;
-    bindGroupLayoutEntry.visibility = WGPUShaderStage_Fragment;
+    bindGroupLayoutEntry.visibility = WGPUShaderStage_Fragment | WGPUShaderStage_Compute;
     bindGroupLayoutEntry.texture.nextInChain = nullptr;
     bindGroupLayoutEntry.texture.sampleType = WGPUTextureSampleType_Float;
     bindGroupLayoutEntry.texture.viewDimension = WGPUTextureViewDimension_2D;
@@ -207,6 +261,20 @@ WGPUBindGroupLayoutEntry WgBindGroup::makeBindGroupLayoutEntryTextureView(uint32
 }
 
 
+WGPUBindGroupLayoutEntry WgBindGroup::makeBindGroupLayoutEntryStorageTexture(uint32_t binding)
+{
+    WGPUBindGroupLayoutEntry bindGroupLayoutEntry{};
+    bindGroupLayoutEntry.nextInChain = nullptr;
+    bindGroupLayoutEntry.binding = binding;
+    bindGroupLayoutEntry.visibility = WGPUShaderStage_Fragment | WGPUShaderStage_Compute;
+    bindGroupLayoutEntry.storageTexture.nextInChain = nullptr;
+    bindGroupLayoutEntry.storageTexture.access = WGPUStorageTextureAccess_ReadWrite;
+    bindGroupLayoutEntry.storageTexture.format = WGPUTextureFormat_RGBA8Unorm;
+    bindGroupLayoutEntry.storageTexture.viewDimension = WGPUTextureViewDimension_2D;
+    return bindGroupLayoutEntry;
+}
+
+
 WGPUBuffer WgBindGroup::createBuffer(WGPUDevice device, WGPUQueue queue, const void *data, size_t size)
 {
     WGPUBufferDescriptor bufferDescriptor{};
@@ -226,7 +294,7 @@ WGPUBindGroup WgBindGroup::createBindGroup(WGPUDevice device, WGPUBindGroupLayou
 {
     WGPUBindGroupDescriptor bindGroupDesc{};
     bindGroupDesc.nextInChain = nullptr;
-    bindGroupDesc.label = "The binding group sampler";
+    bindGroupDesc.label = "The binding group";
     bindGroupDesc.layout = layout;
     bindGroupDesc.entryCount = count;
     bindGroupDesc.entries = bindGroupEntries;
@@ -283,7 +351,7 @@ WGPUPipelineLayout WgPipeline::createPipelineLayout(WGPUDevice device, const WGP
 {
     WGPUPipelineLayoutDescriptor pipelineLayoutDesc{};
     pipelineLayoutDesc.nextInChain = nullptr;
-    pipelineLayoutDesc.label = "The Pipeline layout";
+    pipelineLayoutDesc.label = "The pipeline layout";
     pipelineLayoutDesc.bindGroupLayoutCount = count;
     pipelineLayoutDesc.bindGroupLayouts = bindGroupLayouts;
     return wgpuDeviceCreatePipelineLayout(device, &pipelineLayoutDesc);
@@ -322,11 +390,11 @@ void WgPipeline::destroyShaderModule(WGPUShaderModule& shaderModule)
 // render pipeline
 //*****************************************************************************
 
-void WgRenderPipeline::allocate(WGPUDevice device, 
+void WgRenderPipeline::allocate(WGPUDevice device,
-                          WGPUVertexBufferLayout vertexBufferLayouts[], uint32_t attribsCount,
+                                WGPUVertexBufferLayout vertexBufferLayouts[], uint32_t attribsCount,
-                          WGPUBindGroupLayout bindGroupLayouts[], uint32_t bindGroupsCount,
+                                WGPUBindGroupLayout bindGroupLayouts[], uint32_t bindGroupsCount,
-                          WGPUCompareFunction stencilCompareFunction, WGPUStencilOperation stencilOperation,
+                                WGPUCompareFunction stencilCompareFunction, WGPUStencilOperation stencilOperation,
-                          const char* shaderSource, const char* shaderLabel, const char* pipelineLabel)
+                                const char* shaderSource, const char* shaderLabel, const char* pipelineLabel)
 {
     mShaderModule = createShaderModule(device, shaderSource, shaderLabel);
     assert(mShaderModule);
@@ -372,7 +440,8 @@ WGPUColorTargetState WgRenderPipeline::makeColorTargetState(const WGPUBlendState
 {
     WGPUColorTargetState colorTargetState{};
     colorTargetState.nextInChain = nullptr;
-    colorTargetState.format = WGPUTextureFormat_BGRA8Unorm; // (WGPUTextureFormat_BGRA8UnormSrgb)
+    //colorTargetState.format = WGPUTextureFormat_BGRA8Unorm; // (WGPUTextureFormat_BGRA8UnormSrgb)
+    colorTargetState.format = WGPUTextureFormat_RGBA8Unorm; // (WGPUTextureFormat_BGRA8UnormSrgb)
     colorTargetState.blend = blendState;
     colorTargetState.writeMask = WGPUColorWriteMask_All;
     return colorTargetState;
@@ -415,6 +484,7 @@ WGPUPrimitiveState WgRenderPipeline::makePrimitiveState()
     return primitiveState;
 }
 
+
 WGPUDepthStencilState WgRenderPipeline::makeDepthStencilState(WGPUCompareFunction compare, WGPUStencilOperation operation)
 {
     WGPUDepthStencilState depthStencilState{};
@@ -463,11 +533,12 @@ WGPUFragmentState WgRenderPipeline::makeFragmentState(WGPUShaderModule shaderMod
     return fragmentState;
 }
 
+
 WGPURenderPipeline WgRenderPipeline::createRenderPipeline(WGPUDevice device,
-                                                    WGPUVertexBufferLayout vertexBufferLayouts[], uint32_t attribsCount,
+                                                          WGPUVertexBufferLayout vertexBufferLayouts[], uint32_t attribsCount,
-                                                    WGPUCompareFunction stencilCompareFunction, WGPUStencilOperation stencilOperation,
+                                                          WGPUCompareFunction stencilCompareFunction, WGPUStencilOperation stencilOperation,
-                                                    WGPUPipelineLayout pipelineLayout, WGPUShaderModule shaderModule,
+                                                          WGPUPipelineLayout pipelineLayout, WGPUShaderModule shaderModule,
-                                                    const char* pipelineName)
+                                                          const char* pipelineName)
 {
     WGPUBlendState blendState = makeBlendState();
     WGPUColorTargetState colorTargetStates[] = { 
@@ -492,8 +563,50 @@ WGPURenderPipeline WgRenderPipeline::createRenderPipeline(WGPUDevice device,
     return wgpuDeviceCreateRenderPipeline(device, &renderPipelineDesc);
 }
 
+
 void WgRenderPipeline::destroyRenderPipeline(WGPURenderPipeline& renderPipeline)
 {
     if (renderPipeline) wgpuRenderPipelineRelease(renderPipeline);
     renderPipeline = nullptr;
-}
+}
+
+//*****************************************************************************
+// 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);
+}

src/renderer/wg_engine/tvgWgCommon.h

@@ -46,6 +46,11 @@ struct WgContext {
     void release();
 
     void executeCommandEncoder(WGPUCommandEncoder commandEncoder);
+
+    WGPUTexture createTexture2d(WGPUTextureUsage usage, WGPUTextureFormat format, uint32_t width, uint32_t height, char const * label);
+    WGPUTextureView createTextureView2d(WGPUTexture texture, WGPU_NULLABLE char const * label);
+    void releaseTexture(WGPUTexture& texture);
+    void releaseTextureView(WGPUTextureView& textureView);
 };
 
 struct WgBindGroup
@@ -53,6 +58,7 @@ struct WgBindGroup
     WGPUBindGroup mBindGroup{};
 
     void set(WGPURenderPassEncoder encoder, uint32_t groupIndex);
+    void set(WGPUComputePassEncoder encoder, uint32_t groupIndex);
 
     static WGPUBindGroupEntry makeBindGroupEntryBuffer(uint32_t binding, WGPUBuffer buffer);
     static WGPUBindGroupEntry makeBindGroupEntrySampler(uint32_t binding, WGPUSampler sampler);
@@ -60,7 +66,8 @@ struct WgBindGroup
 
     static WGPUBindGroupLayoutEntry makeBindGroupLayoutEntryBuffer(uint32_t binding);
     static WGPUBindGroupLayoutEntry makeBindGroupLayoutEntrySampler(uint32_t binding);
-    static WGPUBindGroupLayoutEntry makeBindGroupLayoutEntryTextureView(uint32_t binding);
+    static WGPUBindGroupLayoutEntry makeBindGroupLayoutEntryTexture(uint32_t binding);
+    static WGPUBindGroupLayoutEntry makeBindGroupLayoutEntryStorageTexture(uint32_t binding);
 
     static WGPUBuffer createBuffer(WGPUDevice device, WGPUQueue queue, const void *data, size_t size);
     static WGPUBindGroup createBindGroup(WGPUDevice device, WGPUBindGroupLayout layout, const WGPUBindGroupEntry* bindGroupEntries, uint32_t count);
@@ -116,4 +123,17 @@ public:
     static void destroyRenderPipeline(WGPURenderPipeline& renderPipeline);
 };
 
+struct WgComputePipeline: public WgPipeline
+{
+protected:
+    WGPUComputePipeline mComputePipeline{};
+    void allocate(WGPUDevice device,
+                  WGPUBindGroupLayout bindGroupLayouts[], uint32_t bindGroupsCount,
+                  const char* shaderSource, const char* shaderLabel, const char* pipelineLabel);
+
+public:
+    void release() override;
+    void set(WGPUComputePassEncoder computePassEncoder);
+};
+
 #endif // _TVG_WG_COMMON_H_

src/renderer/wg_engine/tvgWgPipelines.cpp

@@ -235,7 +235,7 @@ void WgPipelineBlit::initialize(WGPUDevice device)
 
     // bind groups and layouts
     WGPUBindGroupLayout bindGroupLayouts[] = {
-        WgBindGroupBlit::getLayout(device),
+        WgBindGroupTextureSampled::getLayout(device),
         WgBindGroupOpacity::getLayout(device)
     };
 
@@ -269,7 +269,7 @@ void WgPipelineBlitColor::initialize(WGPUDevice device)
 
     // bind groups and layouts
     WGPUBindGroupLayout bindGroupLayouts[] = {
-        WgBindGroupBlit::getLayout(device)
+        WgBindGroupTextureSampled::getLayout(device)
     };
 
     // stencil function
@@ -302,8 +302,8 @@ void WgPipelineComposition::initialize(WGPUDevice device, const char* shaderSrc)
 
     // bind groups and layouts
     WGPUBindGroupLayout bindGroupLayouts[] = {
-        WgBindGroupBlit::getLayout(device),
+        WgBindGroupTextureSampled::getLayout(device),
-        WgBindGroupBlit::getLayout(device)
+        WgBindGroupTextureSampled::getLayout(device)
     };
 
     // stencil function
@@ -323,18 +323,40 @@ void WgPipelineComposition::initialize(WGPUDevice device, const char* shaderSrc)
              shaderSource, shaderLabel, pipelineLabel);
 }
 
+void WgPipelineBlend::initialize(WGPUDevice device, const char* shaderSrc)
+{
+    // bind groups and layouts
+    WGPUBindGroupLayout bindGroupLayouts[] = {
+        //WgBindGroupTexture::getLayout(device),
+        WgBindGroupStorageTexture::getLayout(device),
+        WgBindGroupStorageTexture::getLayout(device)
+    };
+
+    // sheder source and labels
+    auto shaderSource = shaderSrc;
+    auto shaderLabel = "The compute shader blend";
+    auto pipelineLabel = "The compute pipeline blend";
+
+    // allocate all pipeline handles
+    allocate(device,
+             bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts),
+             shaderSource, shaderLabel, pipelineLabel);
+}
+
 //************************************************************************
 // pipelines
 //************************************************************************
 
 void WgPipelines::initialize(WgContext& context)
 {
+    // fill pipelines
     fillShape.initialize(context.device);
     fillStroke.initialize(context.device);
     solid.initialize(context.device);
     linear.initialize(context.device);
     radial.initialize(context.device);
     image.initialize(context.device);
+    // blit pipelines
     blit.initialize(context.device);
     blitColor.initialize(context.device);
     // composition pipelines
@@ -346,6 +368,8 @@ void WgPipelines::initialize(WgContext& context)
     compSubtractMask.initialize(context.device, cShaderSource_PipelineCompSubtractMask);
     compIntersectMask.initialize(context.device, cShaderSource_PipelineCompIntersectMask);
     compDifferenceMask.initialize(context.device, cShaderSource_PipelineCompDifferenceMask);
+    // compute pipelines
+    computeBlend.initialize(context.device, cShaderSource_PipelineComputeBlend);
     // store pipelines to context
     context.pipelines = this;
 }
@@ -353,7 +377,9 @@ void WgPipelines::initialize(WgContext& context)
 
 void WgPipelines::release()
 {
-    WgBindGroupBlit::releaseLayout();
+    WgBindGroupTextureSampled::releaseLayout();
+    WgBindGroupStorageTexture::releaseLayout();
+    WgBindGroupTexture::releaseLayout();
     WgBindGroupOpacity::releaseLayout();
     WgBindGroupPicture::releaseLayout();
     WgBindGroupRadialGradient::releaseLayout();
@@ -361,6 +387,9 @@ void WgPipelines::release()
     WgBindGroupSolidColor::releaseLayout();
     WgBindGroupPaint::releaseLayout();
     WgBindGroupCanvas::releaseLayout();
+    // compute pipelines
+    computeBlend.release();
+    // composition pipelines
     compDifferenceMask.release();
     compIntersectMask.release();
     compSubtractMask.release();
@@ -369,8 +398,10 @@ void WgPipelines::release()
     compLumaMask.release();
     compInvAlphaMask.release();
     compAlphaMask.release();
+    // blit pipelines
     blitColor.release();
     blit.release();
+    // fill pipelines
     image.release();
     radial.release();
     linear.release();

src/renderer/wg_engine/tvgWgPipelines.h

@@ -25,6 +25,10 @@
 
 #include "tvgWgBindGroups.h"
 
+//*****************************************************************************
+// render pipelines
+//*****************************************************************************
+
 struct WgPipelineFillShape: public WgRenderPipeline
 {
     void initialize(WGPUDevice device) override;
@@ -99,11 +103,11 @@ struct WgPipelineBlit: public WgRenderPipeline
 {
     void initialize(WGPUDevice device) override;
     void use(WGPURenderPassEncoder encoder,
-             WgBindGroupBlit& groupBlit,
+             WgBindGroupTextureSampled& groupTexSampled,
              WgBindGroupOpacity& groupOpacity)
     {
         set(encoder);
-        groupBlit.set(encoder, 0);
+        groupTexSampled.set(encoder, 0);
         groupOpacity.set(encoder, 1);
     }
 };
@@ -111,10 +115,11 @@ struct WgPipelineBlit: public WgRenderPipeline
 struct WgPipelineBlitColor: public WgRenderPipeline
 {
     void initialize(WGPUDevice device) override;
-    void use(WGPURenderPassEncoder encoder, WgBindGroupBlit& groupBlit)
+    void use(WGPURenderPassEncoder encoder, 
+             WgBindGroupTextureSampled& groupTexSampled)
     {
         set(encoder);
-        groupBlit.set(encoder, 0);
+        groupTexSampled.set(encoder, 0);
     }
 };
 
@@ -122,14 +127,38 @@ struct WgPipelineComposition: public WgRenderPipeline
 {
     void initialize(WGPUDevice device) override {};
     void initialize(WGPUDevice device, const char* shaderSrc);
-    void use(WGPURenderPassEncoder encoder, WgBindGroupBlit& groupBlitSrc, WgBindGroupBlit& groupBlitMsk)
+    void use(WGPURenderPassEncoder encoder,
+             WgBindGroupTextureSampled& groupTexSampledSrc,
+             WgBindGroupTextureSampled& groupTexSampledMsk)
     {
         set(encoder);
-        groupBlitSrc.set(encoder, 0);
+        groupTexSampledSrc.set(encoder, 0);
-        groupBlitMsk.set(encoder, 1);
+        groupTexSampledMsk.set(encoder, 1);
     }
 };
 
+//*****************************************************************************
+// compute pipelines
+//*****************************************************************************
+
+struct WgPipelineBlend: public WgComputePipeline
+{
+    void initialize(WGPUDevice device) override {};
+    void initialize(WGPUDevice device, const char* shaderSrc);
+    void use(WGPUComputePassEncoder encoder,
+             WgBindGroupStorageTexture& groupTexSrc,
+             WgBindGroupStorageTexture& groupTexDst)
+    {
+        set(encoder);
+        groupTexSrc.set(encoder, 0);
+        groupTexDst.set(encoder, 1);
+    }
+};
+
+//*****************************************************************************
+// pipelines
+//*****************************************************************************
+
 struct WgPipelines
 {
     WgPipelineFillShape fillShape;
@@ -149,6 +178,8 @@ struct WgPipelines
     WgPipelineComposition compSubtractMask;
     WgPipelineComposition compIntersectMask;
     WgPipelineComposition compDifferenceMask;
+    // compute pipelines
+    WgPipelineBlend computeBlend;
 
     void initialize(WgContext& context);
     void release();

src/renderer/wg_engine/tvgWgRenderTarget.cpp

@@ -48,7 +48,7 @@ void WgRenderTarget::initialize(WgContext& context, uint32_t w, uint32_t h)
     textureDescColor.usage = WGPUTextureUsage_RenderAttachment | WGPUTextureUsage_TextureBinding | WGPUTextureUsage_CopyDst | WGPUTextureUsage_StorageBinding;
     textureDescColor.dimension = WGPUTextureDimension_2D;
     textureDescColor.size = { w, h, 1 };
-    textureDescColor.format = WGPUTextureFormat_BGRA8Unorm;
+    textureDescColor.format = WGPUTextureFormat_RGBA8Unorm;
     textureDescColor.mipLevelCount = 1;
     textureDescColor.sampleCount = 1;
     textureDescColor.viewFormatCount = 0;
@@ -59,7 +59,7 @@ void WgRenderTarget::initialize(WgContext& context, uint32_t w, uint32_t h)
     WGPUTextureViewDescriptor textureViewDescColor{};
     textureViewDescColor.nextInChain = nullptr;
     textureViewDescColor.label = "The target texture view color";
-    textureViewDescColor.format = WGPUTextureFormat_BGRA8Unorm;
+    textureViewDescColor.format = WGPUTextureFormat_RGBA8Unorm;
     textureViewDescColor.dimension = WGPUTextureViewDimension_2D;
     textureViewDescColor.baseMipLevel = 0;
     textureViewDescColor.mipLevelCount = 1;
@@ -96,7 +96,9 @@ void WgRenderTarget::initialize(WgContext& context, uint32_t w, uint32_t h)
     textureViewStencil = wgpuTextureCreateView(mTextureStencil, &textureViewDescStencil);
     assert(textureViewStencil);
     // initialize bind group for blitting
-    bindGroupBlit.initialize(context.device, context.queue, sampler, textureViewColor);
+    bindGroupTex.initialize(context.device, context.queue, textureViewColor);
+    bindGroupStorageTex.initialize(context.device, context.queue, textureViewColor);
+    bindGroupTexSampled.initialize(context.device, context.queue, sampler, textureViewColor);
     // initialize window binding groups
     WgShaderTypeMat4x4f viewMat(w, h);
     mBindGroupCanvasWnd.initialize(context.device, context.queue, viewMat);
@@ -111,7 +113,9 @@ void WgRenderTarget::release(WgContext& context)
 {
     mMeshDataCanvasWnd.release(context);
     mBindGroupCanvasWnd.release();
-    bindGroupBlit.release();
+    bindGroupTexSampled.release();
+    bindGroupStorageTex.release();
+    bindGroupTex.release();
     if (mTextureStencil) {
         wgpuTextureDestroy(mTextureStencil);
         wgpuTextureRelease(mTextureStencil);
@@ -167,7 +171,6 @@ void WgRenderTarget::beginRenderPass(WGPUCommandEncoder commandEncoder, WGPUText
     renderPassDesc.depthStencilAttachment = &depthStencilAttachment;
     //renderPassDesc.depthStencilAttachment = nullptr;
     renderPassDesc.occlusionQuerySet = nullptr;
-    renderPassDesc.timestampWriteCount = 0;
     renderPassDesc.timestampWrites = nullptr;
     // begin render pass
     mRenderPassEncoder = wgpuCommandEncoderBeginRenderPass(commandEncoder, &renderPassDesc);
@@ -251,7 +254,7 @@ void WgRenderTarget::renderPicture(WgRenderDataPicture* renderData)
 void WgRenderTarget::blit(WgContext& context, WgRenderTarget* renderTargetSrc, WgBindGroupOpacity* mBindGroupOpacity)
 {
     assert(mRenderPassEncoder);
-    mPipelines->blit.use(mRenderPassEncoder, renderTargetSrc->bindGroupBlit, *mBindGroupOpacity);
+    mPipelines->blit.use(mRenderPassEncoder, renderTargetSrc->bindGroupTexSampled, *mBindGroupOpacity);
     mMeshDataCanvasWnd.drawImage(mRenderPassEncoder);
 }
 
@@ -259,7 +262,7 @@ void WgRenderTarget::blit(WgContext& context, WgRenderTarget* renderTargetSrc, W
 void WgRenderTarget::blitColor(WgContext& context, WgRenderTarget* renderTargetSrc)
 {
     assert(mRenderPassEncoder);
-    mPipelines->blitColor.use(mRenderPassEncoder, renderTargetSrc->bindGroupBlit);
+    mPipelines->blitColor.use(mRenderPassEncoder, renderTargetSrc->bindGroupTexSampled);
     mMeshDataCanvasWnd.drawImage(mRenderPassEncoder);
 }
 
@@ -269,7 +272,7 @@ void WgRenderTarget::compose(WgContext& context, WgRenderTarget* renderTargetSrc
     assert(mRenderPassEncoder);
     WgPipelineComposition* pipeline = mPipelines->getCompositionPipeline(method);
     assert(pipeline);
-    pipeline->use(mRenderPassEncoder, renderTargetSrc->bindGroupBlit, renderTargetMsk->bindGroupBlit);
+    pipeline->use(mRenderPassEncoder, renderTargetSrc->bindGroupTexSampled, renderTargetMsk->bindGroupTexSampled);
     mMeshDataCanvasWnd.drawImage(mRenderPassEncoder);
 }
 

src/renderer/wg_engine/tvgWgRenderTarget.h

@@ -41,7 +41,9 @@ public:
     WGPUSampler sampler{};
     WGPUTextureView textureViewColor{};
     WGPUTextureView textureViewStencil{};
-    WgBindGroupBlit bindGroupBlit;
+    WgBindGroupTexture bindGroupTex;
+    WgBindGroupStorageTexture bindGroupStorageTex;
+    WgBindGroupTextureSampled bindGroupTexSampled;
 public:
     void initialize(WgContext& context, uint32_t w, uint32_t h);
     void release(WgContext& context);

src/renderer/wg_engine/tvgWgRenderer.cpp

@@ -43,6 +43,7 @@ void WgRenderer::initialize()
 {
     mContext.initialize();
     mPipelines.initialize(mContext);
+    mBindGroupOpacityPool.initialize(mContext);
 }
 
 
@@ -50,12 +51,12 @@ void WgRenderer::release()
 {
     mCompositorStack.clear();
     mRenderTargetStack.clear();
-    mBindGroupOpacityPool.release();
+    mBindGroupOpacityPool.release(mContext);
     mRenderTargetPool.release(mContext);
     mRenderTargetRoot.release(mContext);
     mRenderTargetWnd.release(mContext);
-    if (mSwapChain) wgpuSwapChainRelease(mSwapChain);
+    wgpuSurfaceUnconfigure(mSurface);
-    if (mSurface) wgpuSurfaceRelease(mSurface);
+    wgpuSurfaceRelease(mSurface);
     mPipelines.release();
     mContext.release();
 }
@@ -159,7 +160,6 @@ bool WgRenderer::postRender()
 {
     mRenderTargetRoot.endRenderPass();
     mRenderTargetStack.pop();
-    mBindGroupOpacityPool.reset();
     mContext.executeCommandEncoder(mCommandEncoder);
     wgpuCommandEncoderRelease(mCommandEncoder);
     return true;
@@ -211,7 +211,9 @@ bool WgRenderer::clear()
 
 bool WgRenderer::sync()
 {
-    WGPUTextureView backBufferView = wgpuSwapChainGetCurrentTextureView(mSwapChain);
+    WGPUSurfaceTexture backBuffer{};
+    wgpuSurfaceGetCurrentTexture(mSurface, &backBuffer);
+    WGPUTextureView backBufferView = mContext.createTextureView2d(backBuffer.texture, "Surface texture view");
     WGPUCommandEncoderDescriptor commandEncoderDesc{};
     commandEncoderDesc.nextInChain = nullptr;
     commandEncoderDesc.label = "The command encoder";
@@ -222,7 +224,7 @@ bool WgRenderer::sync()
     mContext.executeCommandEncoder(commandEncoder);
     wgpuCommandEncoderRelease(commandEncoder);
     wgpuTextureViewRelease(backBufferView);
-    wgpuSwapChainPresent(mSwapChain);
+    wgpuSurfacePresent(mSurface);
     return true;
 }
 
@@ -260,22 +262,22 @@ bool WgRenderer::target(void* window, uint32_t w, uint32_t h)
     mSurface = wgpuInstanceCreateSurface(mContext.instance, &surfaceDesc);
     assert(mSurface);
 
-    // get preferred format
+    WGPUSurfaceConfiguration surfaceConfiguration{};
-    WGPUTextureFormat swapChainFormat = WGPUTextureFormat_BGRA8Unorm;
+    surfaceConfiguration.nextInChain = nullptr;
-    // swapchain descriptor
+    surfaceConfiguration.device = mContext.device;
-    WGPUSwapChainDescriptor swapChainDesc{};
+    surfaceConfiguration.format = WGPUTextureFormat_RGBA8Unorm;
-    swapChainDesc.nextInChain = nullptr;
+    surfaceConfiguration.usage = WGPUTextureUsage_RenderAttachment;
-    swapChainDesc.label = "The swapchain";
+    surfaceConfiguration.viewFormatCount = 0;
-    swapChainDesc.usage = WGPUTextureUsage_RenderAttachment;
+    surfaceConfiguration.viewFormats = nullptr;
-    swapChainDesc.format = swapChainFormat;
+    surfaceConfiguration.alphaMode = WGPUCompositeAlphaMode_Auto;
-    swapChainDesc.width = mTargetSurface.w;
+    surfaceConfiguration.width = mTargetSurface.w;
-    swapChainDesc.height = mTargetSurface.h;
+    surfaceConfiguration.height = mTargetSurface.h;
-    swapChainDesc.presentMode = WGPUPresentMode_Mailbox;
+    surfaceConfiguration.presentMode = WGPUPresentMode_Mailbox;
-    mSwapChain = wgpuDeviceCreateSwapChain(mContext.device, mSurface, &swapChainDesc);
+    wgpuSurfaceConfigure(mSurface, &surfaceConfiguration);
-    assert(mSwapChain);
 
     mRenderTargetWnd.initialize(mContext, w, h);
     mRenderTargetRoot.initialize(mContext, w, h);
+
     return true;
 }
 

src/renderer/wg_engine/tvgWgRenderer.h

@@ -59,7 +59,7 @@ public:
     static WgRenderer* gen();
     static bool init(uint32_t threads);
     static bool term();
-
+private:
     // render handles
     WGPUCommandEncoder mCommandEncoder{};
     Array<Compositor*> mCompositorStack;
@@ -74,7 +74,6 @@ private:
     WgRenderTarget mRenderTargetRoot;
     WgRenderTarget mRenderTargetWnd;
     WGPUSurface mSurface{};
-    WGPUSwapChain mSwapChain{};
     Surface mTargetSurface;
 };
 

src/renderer/wg_engine/tvgWgShaderSrc.cpp

@@ -707,3 +707,23 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f {
     }
 };
 )";
+
+//************************************************************************
+// cShaderSource_PipelineComputeBlend
+//************************************************************************
+
+// pipeline shader modules blend (simple example)
+const char* cShaderSource_PipelineComputeBlend = R"(
+@group(0) @binding(0) var imageSrc : texture_storage_2d<rgba8unorm, write>;
+@group(1) @binding(0) var imageDst : texture_storage_2d<rgba8unorm, write>;
+
+@compute @workgroup_size(8, 8)
+fn cs_main( @builtin(global_invocation_id) id: vec3u) {
+   let texSize = textureDimensions(imageSrc);
+   if ((id.x >= texSize.x) || (id.y >= texSize.y)) { return; };
+
+   // var src = textureLoad(imageSrc, id.xy);
+   textureStore(imageDst, id.xy, vec4(0.5, 1.0, 0.1, 1.0));
+}
+
+)";

src/renderer/wg_engine/tvgWgShaderSrc.h

@@ -25,26 +25,22 @@
 #ifndef _TVG_WG_SHADER_SRC_H_
 #define _TVG_WG_SHADER_SRC_H_
 
-// pipeline shader module fill
+//*****************************************************************************
+// render shader modules
+//*****************************************************************************
+
+// pipeline shader modules fill
 extern const char* cShaderSource_PipelineFill;
-
-// pipeline shader module solid
 extern const char* cShaderSource_PipelineSolid;
-
-// pipeline shader module linear
 extern const char* cShaderSource_PipelineLinear;
-
-// pipeline shader module radial
 extern const char* cShaderSource_PipelineRadial;
-
-// pipeline shader module image
 extern const char* cShaderSource_PipelineImage;
 
-// pipeline shader module blit
+// pipeline shader modules blit
 extern const char* cShaderSource_PipelineBlit;
 extern const char* cShaderSource_PipelineBlitColor;
 
-// pipeline shader module composes
+// pipeline shader modules composes
 extern const char* cShaderSource_PipelineCompAlphaMask;
 extern const char* cShaderSource_PipelineCompInvAlphaMask;
 extern const char* cShaderSource_PipelineCompLumaMask;
@@ -54,4 +50,11 @@ extern const char* cShaderSource_PipelineCompSubtractMask;
 extern const char* cShaderSource_PipelineCompIntersectMask;
 extern const char* cShaderSource_PipelineCompDifferenceMask;
 
+//*****************************************************************************
+// compute shader modules
+//*****************************************************************************
+
+// pipeline shader modules blend
+extern const char* cShaderSource_PipelineComputeBlend;
+
 #endif // _TVG_WG_SHADER_SRC_H_