wg_engine: antialiasing

[issues 1479: antialiasing](#1479)

Anti-aliasing implementation
Implements antialiasing as a post process on cimpute shaders and original render target with scale of 2x.
Can be modified as an external settings

src/renderer/wg_engine/tvgWgCommon.cpp

@@ -158,6 +158,24 @@ WGPUTexture WgContext::createTexture2d(WGPUTextureUsageFlags usage, WGPUTextureF
 }
 
 
+WGPUTexture WgContext::createTexture2dMS(WGPUTextureUsageFlags usage, WGPUTextureFormat format, uint32_t width, uint32_t height, uint32_t sc, 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 = sc;
+    textureDesc.viewFormatCount = 0;
+    textureDesc.viewFormats = nullptr;
+    return wgpuDeviceCreateTexture(device, &textureDesc);
+}
+
+
 WGPUTextureView WgContext::createTextureView2d(WGPUTexture texture, char const * label)
 {
     WGPUTextureViewDescriptor textureViewDescColor{};

src/renderer/wg_engine/tvgWgCommon.h

@@ -52,6 +52,7 @@ struct WgContext {
 
     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);
 

src/renderer/wg_engine/tvgWgPipelines.cpp

@@ -291,6 +291,26 @@ void WgPipelineCompose::initialize(WGPUDevice device)
              shaderSource, shaderLabel, pipelineLabel);
 }
 
+
+void WgPipelineAntiAliasing::initialize(WGPUDevice device)
+{
+    // bind groups and layouts
+    WGPUBindGroupLayout bindGroupLayouts[] = {
+        WgBindGroupTextureStorage::getLayout(device),
+        WgBindGroupTextureStorage::getLayout(device)
+    };
+
+    // sheder 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
 //************************************************************************
@@ -308,6 +328,7 @@ void WgPipelines::initialize(WgContext& context)
     computeClear.initialize(context.device);
     computeBlend.initialize(context.device);
     computeCompose.initialize(context.device);
+    computeAntiAliasing.initialize(context.device);
     // store pipelines to context
     context.pipelines = this;
 }
@@ -326,6 +347,7 @@ void WgPipelines::release()
     WgBindGroupPaint::releaseLayout();
     WgBindGroupCanvas::releaseLayout();
     // compute pipelines
+    computeAntiAliasing.release();
     computeCompose.release();
     computeBlend.release();
     computeClear.release();

src/renderer/wg_engine/tvgWgPipelines.h

@@ -106,8 +106,7 @@ struct WgPipelineImage: public WgRenderPipeline
 struct WgPipelineClear: public WgComputePipeline
 {
     void initialize(WGPUDevice device) override;
-    void use(WGPUComputePassEncoder encoder,
-             WgBindGroupTextureStorage& groupTexDst)
+    void use(WGPUComputePassEncoder encoder, WgBindGroupTextureStorage& groupTexDst)
     {
         set(encoder);
         groupTexDst.set(encoder, 0);
@@ -118,10 +117,7 @@ struct WgPipelineClear: public WgComputePipeline
 struct WgPipelineBlend: public WgComputePipeline
 {
     void initialize(WGPUDevice device) override;
-    void use(WGPUComputePassEncoder encoder,
-             WgBindGroupTextureStorage& groupTexSrc,
-             WgBindGroupTextureStorage& groupTexDst,
-             WgBindGroupBlendMethod& blendMethod)
+    void use(WGPUComputePassEncoder encoder, WgBindGroupTextureStorage& groupTexSrc, WgBindGroupTextureStorage& groupTexDst, WgBindGroupBlendMethod& blendMethod)
     {
         set(encoder);
         groupTexSrc.set(encoder, 0);
@@ -134,11 +130,7 @@ struct WgPipelineBlend: public WgComputePipeline
 struct WgPipelineCompose: public WgComputePipeline
 {
     void initialize(WGPUDevice device) override;
-    void use(WGPUComputePassEncoder encoder,
-             WgBindGroupTextureStorage& groupTexSrc,
-             WgBindGroupTextureStorage& groupTexMsk,
-             WgBindGroupCompositeMethod& groupComposeMethod,
-             WgBindGroupOpacity& groupOpacity)
+    void use(WGPUComputePassEncoder encoder, WgBindGroupTextureStorage& groupTexSrc, WgBindGroupTextureStorage& groupTexMsk, WgBindGroupCompositeMethod& groupComposeMethod, WgBindGroupOpacity& groupOpacity)
     {
         set(encoder);
         groupTexSrc.set(encoder, 0);
@@ -148,6 +140,18 @@ struct WgPipelineCompose: public WgComputePipeline
     }
 };
 
+
+struct WgPipelineAntiAliasing: public WgComputePipeline
+{
+    void initialize(WGPUDevice device) override;
+    void use(WGPUComputePassEncoder encoder, WgBindGroupTextureStorage& groupTexSrc, WgBindGroupTextureStorage& groupTexDst)
+    {
+        set(encoder);
+        groupTexSrc.set(encoder, 0);
+        groupTexDst.set(encoder, 1);
+    }
+};
+
 //*****************************************************************************
 // pipelines
 //*****************************************************************************
@@ -165,6 +169,7 @@ struct WgPipelines
     WgPipelineClear computeClear;
     WgPipelineBlend computeBlend;
     WgPipelineCompose computeCompose;
+    WgPipelineAntiAliasing computeAntiAliasing;
 
     void initialize(WgContext& context);
     void release();

src/renderer/wg_engine/tvgWgRenderTarget.cpp

@@ -26,18 +26,18 @@
 // render target
 //*****************************************************************************
 
-void WgRenderTarget::initialize(WgContext& context, uint32_t w, uint32_t h)
+void WgRenderTarget::initialize(WgContext& context, uint32_t w, uint32_t h, uint32_t samples)
 {
     release(context);
     // create color and stencil textures
     texColor = context.createTexture2d(
         WGPUTextureUsage_RenderAttachment | WGPUTextureUsage_StorageBinding,
         WGPUTextureFormat_RGBA8Unorm, 
-        w, h, "The target texture color");
+        w * samples, h * samples, "The target texture color");
     texStencil = context.createTexture2d(
         WGPUTextureUsage_RenderAttachment,
         WGPUTextureFormat_Stencil8, 
-        w, h, "The target texture stencil");
+        w * samples, h * samples, "The target texture stencil");
     assert(texColor);
     assert(texStencil);
     texViewColor = context.createTextureView2d(texColor, "The target texture view color");
@@ -183,19 +183,19 @@ void WgRenderTarget::endRenderPass(WGPURenderPassEncoder renderPassEncoder)
 // render storage
 //*****************************************************************************
 
- void WgRenderStorage::initialize(WgContext& context, uint32_t w, uint32_t h)
+ void WgRenderStorage::initialize(WgContext& context, uint32_t w, uint32_t h, uint32_t samples)
  {
     release(context);
     // store target storage size
-    width = w;
-    height = h;
+    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
     texStorage = context.createTexture2d(
-        WGPUTextureUsage_StorageBinding | WGPUTextureUsage_CopySrc,
+        WGPUTextureUsage_StorageBinding | WGPUTextureUsage_CopySrc | WGPUTextureUsage_CopyDst,
         WGPUTextureFormat_RGBA8Unorm,
-        w, h, "The target texture storage color");
+        width, height, "The target texture storage color");
     assert(texStorage);
     texViewStorage = context.createTextureView2d(texStorage, "The target texture storage view color");
     assert(texViewStorage);
@@ -258,6 +258,17 @@ void WgRenderStorage::compose(WGPUCommandEncoder commandEncoder, WgRenderStorage
 };
 
 
+void WgRenderStorage::antialias(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc)
+{
+    assert(commandEncoder);
+    assert(targetSrc);
+    WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder);
+    mPipelines->computeAntiAliasing.use(computePassEncoder, targetSrc->bindGroupTexStorage, bindGroupTexStorage);
+    dispatchWorkgroups(computePassEncoder);
+    endRenderPass(computePassEncoder);
+}
+
+
 void WgRenderStorage::dispatchWorkgroups(WGPUComputePassEncoder computePassEncoder)
 {
     assert(computePassEncoder);

src/renderer/wg_engine/tvgWgRenderTarget.h

@@ -37,7 +37,7 @@ public:
     WGPUTextureView texViewStencil{};
     WgBindGroupTextureStorage bindGroupTexStorage;
 public:
-    void initialize(WgContext& context, uint32_t w, uint32_t h);
+    void initialize(WgContext& context, uint32_t w, uint32_t h, uint32_t samples = 1);
     void release(WgContext& context);
 
     void renderShape(WGPUCommandEncoder commandEncoder, WgRenderDataShape* renderData);
@@ -64,13 +64,14 @@ public:
     uint32_t workgroupsCountX{};
     uint32_t workgroupsCountY{};
 public:
-    void initialize(WgContext& context, uint32_t w, uint32_t h);
+    void initialize(WgContext& context, uint32_t w, uint32_t h, uint32_t samples = 1);
     void release(WgContext& context);
 
     void clear(WGPUCommandEncoder commandEncoder);
     void blend(WGPUCommandEncoder commandEncoder, WgRenderTarget* targetSrc, WgBindGroupBlendMethod* blendMethod);
     void blend(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc, WgBindGroupBlendMethod* blendMethod);
     void compose(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetMsk, WgBindGroupCompositeMethod* composeMethod, WgBindGroupOpacity* opacity);
+    void antialias(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc);
 private:
     void dispatchWorkgroups(WGPUComputePassEncoder computePassEncoder);
 

src/renderer/wg_engine/tvgWgRenderer.cpp

@@ -27,6 +27,8 @@
 #include <windows.h>
 #endif
 
+#define WG_SSAA_SAMPLES (2)
+
 WgRenderer::WgRenderer()
 {
     initialize();
@@ -58,6 +60,7 @@ void WgRenderer::release()
     mBlendMethodPool.release(mContext);
     mOpacityPool.release(mContext);
     mRenderStorageRoot.release(mContext);
+    mRenderStorageScreen.release(mContext);
     mRenderTarget.release(mContext);
     wgpuSurfaceUnconfigure(mSurface);
     wgpuSurfaceRelease(mSurface);
@@ -227,8 +230,10 @@ bool WgRenderer::sync()
     commandEncoderDesc.label = "The command encoder";
     WGPUCommandEncoder commandEncoder = wgpuDeviceCreateCommandEncoder(mContext.device, &commandEncoderDesc);
 
+    mRenderStorageScreen.antialias(commandEncoder, &mRenderStorageRoot);
+
     WGPUImageCopyTexture source{};
-    source.texture = mRenderStorageRoot.texStorage;
+    source.texture = mRenderStorageScreen.texStorage;
     WGPUImageCopyTexture dest{};
     dest.texture = backBuffer.texture;
     WGPUExtent3D copySize{};
@@ -252,7 +257,7 @@ bool WgRenderer::target(uint32_t* buffer, uint32_t stride, uint32_t w, uint32_t
     mTargetSurface.w = w;
     mTargetSurface.h = h;
 
-    mRenderTarget.initialize(mContext, w, h);
+    mRenderTarget.initialize(mContext, w * WG_SSAA_SAMPLES, h * WG_SSAA_SAMPLES);
     return true;
 }
 
@@ -291,8 +296,9 @@ bool WgRenderer::target(void* window, uint32_t w, uint32_t h)
     surfaceConfiguration.presentMode = WGPUPresentMode_Mailbox;
     wgpuSurfaceConfigure(mSurface, &surfaceConfiguration);
 
-    mRenderTarget.initialize(mContext, w, h);
-    mRenderStorageRoot.initialize(mContext, w, h);
+    mRenderTarget.initialize(mContext, w, h, WG_SSAA_SAMPLES);
+    mRenderStorageRoot.initialize(mContext, w, h, WG_SSAA_SAMPLES);
+    mRenderStorageScreen.initialize(mContext, w, h);
 
     return true;
 }
@@ -312,7 +318,7 @@ bool WgRenderer::beginComposite(TVG_UNUSED Compositor* cmp, TVG_UNUSED Composite
     cmp->opacity = opacity;
 
     // allocate new render storage and push it to top of render tree
-    WgRenderStorage* renderStorage = mRenderStoragePool.allocate(mContext, mTargetSurface.w, mTargetSurface.h);
+    WgRenderStorage* renderStorage = mRenderStoragePool.allocate(mContext, mTargetSurface.w * WG_SSAA_SAMPLES, mTargetSurface.h * WG_SSAA_SAMPLES);
     renderStorage->clear(mCommandEncoder);
     mRenderStorageStack.push(renderStorage);
 

src/renderer/wg_engine/tvgWgRenderer.h

@@ -64,6 +64,7 @@ private:
     WGPUCommandEncoder mCommandEncoder{};
     WgRenderTarget mRenderTarget;
     WgRenderStorage mRenderStorageRoot;
+    WgRenderStorage mRenderStorageScreen;
     WgRenderStoragePool mRenderStoragePool;
     WgBindGroupOpacityPool mOpacityPool;
     WgBindGroupBlendMethodPool mBlendMethodPool;

src/renderer/wg_engine/tvgWgShaderSrc.cpp

@@ -428,3 +428,26 @@ fn cs_main( @builtin(global_invocation_id) id: vec3u) {
     textureStore(imageSrc, id.xy, vec4f(color, alpha * opacity));
 }
 )";
+
+// pipeline shader modules anti-aliasing
+const char* cShaderSource_PipelineComputeAntiAlias = R"(
+@group(0) @binding(0) var imageSrc : texture_storage_2d<rgba8unorm, read_write>;
+@group(1) @binding(0) var imageDst : texture_storage_2d<rgba8unorm, read_write>;
+
+@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));
+}
+)";

src/renderer/wg_engine/tvgWgShaderSrc.h

@@ -44,5 +44,6 @@ extern const char* cShaderSource_PipelineImage;
 extern const char* cShaderSource_PipelineComputeClear;
 extern const char* cShaderSource_PipelineComputeBlend;
 extern const char* cShaderSource_PipelineComputeCompose;
+extern const char* cShaderSource_PipelineComputeAntiAlias;
 
 #endif // _TVG_WG_SHADER_SRC_H_