diff --git a/src/renderer/wg_engine/tvgWgCommon.cpp b/src/renderer/wg_engine/tvgWgCommon.cpp index ae4e5faa..a074276b 100644 --- a/src/renderer/wg_engine/tvgWgCommon.cpp +++ b/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{}; diff --git a/src/renderer/wg_engine/tvgWgCommon.h b/src/renderer/wg_engine/tvgWgCommon.h index a8cfbe39..1fbc9191 100644 --- a/src/renderer/wg_engine/tvgWgCommon.h +++ b/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); diff --git a/src/renderer/wg_engine/tvgWgPipelines.cpp b/src/renderer/wg_engine/tvgWgPipelines.cpp index 209a3692..77af9547 100644 --- a/src/renderer/wg_engine/tvgWgPipelines.cpp +++ b/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(); diff --git a/src/renderer/wg_engine/tvgWgPipelines.h b/src/renderer/wg_engine/tvgWgPipelines.h index 99097828..5ec4fe2c 100644 --- a/src/renderer/wg_engine/tvgWgPipelines.h +++ b/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(); diff --git a/src/renderer/wg_engine/tvgWgRenderTarget.cpp b/src/renderer/wg_engine/tvgWgRenderTarget.cpp index 11c9287f..1cc72913 100644 --- a/src/renderer/wg_engine/tvgWgRenderTarget.cpp +++ b/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); diff --git a/src/renderer/wg_engine/tvgWgRenderTarget.h b/src/renderer/wg_engine/tvgWgRenderTarget.h index a4c85b53..35f5b049 100644 --- a/src/renderer/wg_engine/tvgWgRenderTarget.h +++ b/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); diff --git a/src/renderer/wg_engine/tvgWgRenderer.cpp b/src/renderer/wg_engine/tvgWgRenderer.cpp index 19cc1d70..20bfbfcc 100644 --- a/src/renderer/wg_engine/tvgWgRenderer.cpp +++ b/src/renderer/wg_engine/tvgWgRenderer.cpp @@ -27,6 +27,8 @@ #include #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); diff --git a/src/renderer/wg_engine/tvgWgRenderer.h b/src/renderer/wg_engine/tvgWgRenderer.h index 3a4abbf1..1a6d9e84 100644 --- a/src/renderer/wg_engine/tvgWgRenderer.h +++ b/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; diff --git a/src/renderer/wg_engine/tvgWgShaderSrc.cpp b/src/renderer/wg_engine/tvgWgShaderSrc.cpp index 2f10b5e3..d1fd4ba8 100644 --- a/src/renderer/wg_engine/tvgWgShaderSrc.cpp +++ b/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; +@group(1) @binding(0) var imageDst : texture_storage_2d; + +@compute @workgroup_size(8, 8) +fn cs_main( @builtin(global_invocation_id) id: vec3u) { + let texSizeSrc = textureDimensions(imageSrc); + let texSizeDst = textureDimensions(imageDst); + if ((id.x >= texSizeDst.x) || (id.y >= texSizeDst.y)) { return; }; + + let samples = u32(texSizeSrc.x / texSizeDst.x); + var color = vec4f(0); + for (var i = 0u; i < samples; i++) { + for (var j = 0u; j < samples; j++) { + color += textureLoad(imageSrc, vec2(id.x * samples + j, id.y * samples + i)); + } + } + + textureStore(imageDst, id.xy, color / f32(samples * samples)); +} +)"; diff --git a/src/renderer/wg_engine/tvgWgShaderSrc.h b/src/renderer/wg_engine/tvgWgShaderSrc.h index 8f9476f1..7eec4b62 100644 --- a/src/renderer/wg_engine/tvgWgShaderSrc.h +++ b/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_