mirror of
https://github.com/thorvg/thorvg.git
synced 2025-06-14 12:04:29 +00:00
wg_engine: gaussian blur basic implementation
Introduce blur effect for webgpu renderer Issue: https://github.com/thorvg/thorvg/issues/3054
This commit is contained in:
Sergii Liebodkin
Hermet Park
parent
04530ab6c2
commit
857f1404e1
12 changed files with 474 additions and 34 deletions
|
|
@ -22,6 +22,7 @@
|
|||
|
||||
#include "tvgWgCompositor.h"
|
||||
#include "tvgWgShaderTypes.h"
|
||||
#include <iostream>
|
||||
|
||||
void WgCompositor::initialize(WgContext& context, uint32_t width, uint32_t height)
|
||||
{
|
||||
|
|
@ -115,8 +116,8 @@ RenderRegion WgCompositor::shrinkRenderRegion(RenderRegion& rect)
|
|||
// cut viewport to screen dimensions
|
||||
int32_t xmin = std::max(0, std::min((int32_t)width, rect.x));
|
||||
int32_t ymin = std::max(0, std::min((int32_t)height, rect.y));
|
||||
int32_t xmax = std::max(0, std::min((int32_t)width, rect.x + rect.w));
|
||||
int32_t ymax = std::max(0, std::min((int32_t)height, rect.y + rect.h));
|
||||
int32_t xmax = std::max(xmin, std::min((int32_t)width, rect.x + rect.w));
|
||||
int32_t ymax = std::max(ymin, std::min((int32_t)height, rect.y + rect.h));
|
||||
return { xmin, ymin, xmax - xmin, ymax - ymin };
|
||||
}
|
||||
|
||||
|
|
@ -150,11 +151,13 @@ void WgCompositor::beginRenderPass(WGPUCommandEncoder commandEncoder, WgRenderSt
|
|||
|
||||
void WgCompositor::endRenderPass()
|
||||
{
|
||||
assert(renderPassEncoder);
|
||||
wgpuRenderPassEncoderEnd(renderPassEncoder);
|
||||
wgpuRenderPassEncoderRelease(renderPassEncoder);
|
||||
this->renderPassEncoder = nullptr;
|
||||
this->currentTarget = nullptr;
|
||||
if (currentTarget) {
|
||||
assert(renderPassEncoder);
|
||||
wgpuRenderPassEncoderEnd(renderPassEncoder);
|
||||
wgpuRenderPassEncoderRelease(renderPassEncoder);
|
||||
this->renderPassEncoder = nullptr;
|
||||
this->currentTarget = nullptr;
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
|
@ -733,4 +736,52 @@ void WgCompositor::clearClipPath(WgContext& context, WgRenderDataPaint* paint)
|
|||
wgpuRenderPassEncoderSetPipeline(renderPassEncoder, pipelines.clear_depth);
|
||||
renderData->meshDataBBox.drawFan(context, renderPassEncoder);
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
void WgCompositor::gaussianBlur(WgContext& context, WgRenderStorage* dst, const RenderEffectGaussianBlur* params, const WgCompose* compose)
|
||||
{
|
||||
assert(dst);
|
||||
assert(params);
|
||||
assert(params->rd);
|
||||
assert(compose->rdViewport);
|
||||
assert(!renderPassEncoder);
|
||||
auto renderDataGaussian = (WgRenderDataGaussian*)params->rd;
|
||||
auto viewport = compose->rdViewport;
|
||||
for (uint32_t level = 0; level < renderDataGaussian->level; level++) {
|
||||
// horizontal blur
|
||||
if (params->direction != 2) {
|
||||
const WGPUImageCopyTexture texSrc { .texture = dst->texture };
|
||||
const WGPUImageCopyTexture texDst { .texture = storageDstCopy.texture };
|
||||
const WGPUExtent3D copySize { .width = width, .height = height, .depthOrArrayLayers = 1 };
|
||||
wgpuCommandEncoderCopyTextureToTexture(commandEncoder, &texSrc, &texDst, ©Size);
|
||||
WGPUComputePassDescriptor computePassDesc{ .label = "Compute pass gaussian blur horizontal" };
|
||||
WGPUComputePassEncoder computePassEncoder = wgpuCommandEncoderBeginComputePass(commandEncoder, &computePassDesc);
|
||||
wgpuComputePassEncoderSetBindGroup(computePassEncoder, 0, storageDstCopy.bindGroupRead, 0, nullptr);
|
||||
wgpuComputePassEncoderSetBindGroup(computePassEncoder, 1, dst->bindGroupWrite, 0, nullptr);
|
||||
wgpuComputePassEncoderSetBindGroup(computePassEncoder, 2, renderDataGaussian->bindGroupGaussian, 0, nullptr);
|
||||
wgpuComputePassEncoderSetBindGroup(computePassEncoder, 3, viewport->bindGroupViewport, 0, nullptr);
|
||||
wgpuComputePassEncoderSetPipeline(computePassEncoder, pipelines.gaussian_horz);
|
||||
wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, width / 16, height / 16, 1);
|
||||
wgpuComputePassEncoderEnd(computePassEncoder);
|
||||
wgpuComputePassEncoderRelease(computePassEncoder);
|
||||
}
|
||||
// vertical blur
|
||||
if (params->direction != 1) {
|
||||
const WGPUImageCopyTexture texSrc { .texture = dst->texture };
|
||||
const WGPUImageCopyTexture texDst { .texture = storageDstCopy.texture };
|
||||
const WGPUExtent3D copySize { .width = width, .height = height, .depthOrArrayLayers = 1 };
|
||||
wgpuCommandEncoderCopyTextureToTexture(commandEncoder, &texSrc, &texDst, ©Size);
|
||||
WGPUComputePassDescriptor computePassDesc{ .label = "Compute pass gaussian blur vertical" };
|
||||
WGPUComputePassEncoder computePassEncoder = wgpuCommandEncoderBeginComputePass(commandEncoder, &computePassDesc);
|
||||
wgpuComputePassEncoderSetBindGroup(computePassEncoder, 0, storageDstCopy.bindGroupRead, 0, nullptr);
|
||||
wgpuComputePassEncoderSetBindGroup(computePassEncoder, 1, dst->bindGroupWrite, 0, nullptr);
|
||||
wgpuComputePassEncoderSetBindGroup(computePassEncoder, 2, renderDataGaussian->bindGroupGaussian, 0, nullptr);
|
||||
wgpuComputePassEncoderSetBindGroup(computePassEncoder, 3, viewport->bindGroupViewport, 0, nullptr);
|
||||
wgpuComputePassEncoderSetPipeline(computePassEncoder, pipelines.gaussian_vert);
|
||||
wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, width / 16, height / 16, 1);
|
||||
wgpuComputePassEncoderEnd(computePassEncoder);
|
||||
wgpuComputePassEncoderRelease(computePassEncoder);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
|
@ -30,6 +30,7 @@ struct WgCompose: RenderCompositor
|
|||
{
|
||||
BlendMethod blend{};
|
||||
RenderRegion aabb{};
|
||||
WgRenderDataViewport* rdViewport;
|
||||
};
|
||||
|
||||
class WgCompositor
|
||||
|
|
@ -104,6 +105,9 @@ public:
|
|||
|
||||
// blit render storage to texture view (f.e. screen buffer)
|
||||
void blit(WgContext& context, WGPUCommandEncoder encoder, WgRenderStorage* src, WGPUTextureView dstView);
|
||||
|
||||
// effects
|
||||
void gaussianBlur(WgContext& context, WgRenderStorage* dst, const RenderEffectGaussianBlur* params, const WgCompose* compose);
|
||||
};
|
||||
|
||||
#endif // _TVG_WG_COMPOSITOR_H_
|
||||
|
|
|
|||
|
|
@ -183,6 +183,8 @@ void WgPipelines::initialize(WgContext& context)
|
|||
const WGPUBindGroupLayout bindGroupLayoutsSceneCompose[] { layouts.layoutTexSampled, layouts.layoutTexSampled };
|
||||
// bind group layouts blit
|
||||
const WGPUBindGroupLayout bindGroupLayoutsBlit[] { layouts.layoutTexSampled };
|
||||
// bind group layouts effects
|
||||
const WGPUBindGroupLayout bindGroupLayoutsGaussian[] { layouts.layoutTexStrorage1RO, layouts.layoutTexStrorage1WO, layouts.layoutBuffer1Un, layouts.layoutBuffer1Un };
|
||||
|
||||
// depth stencil state markup
|
||||
const WGPUDepthStencilState depthStencilStateNonZero = makeDepthStencilState(WGPUCompareFunction_Always, false, WGPUCompareFunction_Always, WGPUStencilOperation_IncrementWrap, WGPUCompareFunction_Always, WGPUStencilOperation_DecrementWrap);
|
||||
|
|
@ -218,6 +220,9 @@ void WgPipelines::initialize(WgContext& context)
|
|||
shader_scene_compose = createShaderModule(context.device, "The shader scene composition", cShaderSrc_Scene_Compose);
|
||||
// shader blit
|
||||
shader_blit = createShaderModule(context.device, "The shader blit", cShaderSrc_Blit);
|
||||
// shader effects
|
||||
shader_gaussian_horz = createShaderModule(context.device, "The shader gaussian horizontal", cShaderSrc_GaussianBlur_Horz);
|
||||
shader_gaussian_vert = createShaderModule(context.device, "The shader gaussian vertical", cShaderSrc_GaussianBlur_Vert);
|
||||
|
||||
// layouts
|
||||
layout_stencil = createPipelineLayout(context.device, bindGroupLayoutsStencil, 2);
|
||||
|
|
@ -236,6 +241,8 @@ void WgPipelines::initialize(WgContext& context)
|
|||
layout_scene_compose = createPipelineLayout(context.device, bindGroupLayoutsSceneCompose, 2);
|
||||
// layout blit
|
||||
layout_blit = createPipelineLayout(context.device, bindGroupLayoutsBlit, 1);
|
||||
// layout effects
|
||||
layout_gaussian = createPipelineLayout(context.device, bindGroupLayoutsGaussian, 4);
|
||||
|
||||
// render pipeline nonzero
|
||||
nonzero = createRenderPipeline(
|
||||
|
|
@ -439,10 +446,21 @@ void WgPipelines::initialize(WgContext& context)
|
|||
layout_blit, vertexBufferLayoutsImage, 2,
|
||||
WGPUColorWriteMask_All, context.preferredFormat, blendStateSrc, // must be preferred screen pixel format
|
||||
depthStencilStateScene, multisampleStateX1);
|
||||
|
||||
// compute pipeline gaussian blur
|
||||
gaussian_horz = createComputePipeline(
|
||||
context.device, "The compute pipeline gaussian blur horizontal",
|
||||
shader_gaussian_horz, "cs_main", layout_gaussian);
|
||||
gaussian_vert = createComputePipeline(
|
||||
context.device, "The compute pipeline gaussian blur vertical",
|
||||
shader_gaussian_vert, "cs_main", layout_gaussian);
|
||||
}
|
||||
|
||||
void WgPipelines::releaseGraphicHandles(WgContext& context)
|
||||
{
|
||||
// pipeline effects
|
||||
releaseComputePipeline(gaussian_vert);
|
||||
releaseComputePipeline(gaussian_horz);
|
||||
// pipeline blit
|
||||
releaseRenderPipeline(blit);
|
||||
// pipelines compose
|
||||
|
|
@ -473,6 +491,7 @@ void WgPipelines::releaseGraphicHandles(WgContext& context)
|
|||
releaseRenderPipeline(evenodd);
|
||||
releaseRenderPipeline(nonzero);
|
||||
// layouts
|
||||
releasePipelineLayout(layout_gaussian);
|
||||
releasePipelineLayout(layout_blit);
|
||||
releasePipelineLayout(layout_scene_compose);
|
||||
releasePipelineLayout(layout_scene_blend);
|
||||
|
|
@ -486,6 +505,8 @@ void WgPipelines::releaseGraphicHandles(WgContext& context)
|
|||
releasePipelineLayout(layout_depth);
|
||||
releasePipelineLayout(layout_stencil);
|
||||
// shaders
|
||||
releaseShaderModule(shader_gaussian_horz);
|
||||
releaseShaderModule(shader_gaussian_vert);
|
||||
releaseShaderModule(shader_blit);
|
||||
releaseShaderModule(shader_scene_compose);
|
||||
releaseShaderModule(shader_scene_blend);
|
||||
|
|
|
|||
|
|
@ -46,6 +46,9 @@ private:
|
|||
WGPUShaderModule shader_scene_compose{};
|
||||
// shader blit
|
||||
WGPUShaderModule shader_blit{};
|
||||
// shader effects
|
||||
WGPUShaderModule shader_gaussian_horz{};
|
||||
WGPUShaderModule shader_gaussian_vert{};
|
||||
|
||||
// layouts helpers
|
||||
WGPUPipelineLayout layout_stencil{};
|
||||
|
|
@ -64,6 +67,8 @@ private:
|
|||
WGPUPipelineLayout layout_scene_compose{};
|
||||
// layouts blit
|
||||
WGPUPipelineLayout layout_blit{};
|
||||
// layouts effects
|
||||
WGPUPipelineLayout layout_gaussian{};
|
||||
public:
|
||||
// pipelines stencil markup
|
||||
WGPURenderPipeline nonzero{};
|
||||
|
|
@ -91,6 +96,9 @@ public:
|
|||
WGPURenderPipeline scene_compose[11]{};
|
||||
// pipeline blit
|
||||
WGPURenderPipeline blit{};
|
||||
// effects
|
||||
WGPUComputePipeline gaussian_horz{};
|
||||
WGPUComputePipeline gaussian_vert{};
|
||||
private:
|
||||
void releaseGraphicHandles(WgContext& context);
|
||||
WGPUShaderModule createShaderModule(WGPUDevice device, const char* label, const char* code);
|
||||
|
|
|
|||
|
|
@ -369,10 +369,10 @@ void WgRenderDataShape::updateAABB(const Matrix& tr) {
|
|||
Point p1 = Point{pMax.x, pMin.y} * tr;
|
||||
Point p2 = Point{pMin.x, pMax.y} * tr;
|
||||
Point p3 = Point{pMax.x, pMax.y} * tr;
|
||||
aabb.x = std::min({p0.x, p1.x, p2.x, p3.x});
|
||||
aabb.y = std::min({p0.y, p1.y, p2.y, p3.y});
|
||||
aabb.w = std::max({p0.x, p1.x, p2.x, p3.x}) - aabb.x;
|
||||
aabb.h = std::max({p0.y, p1.y, p2.y, p3.y}) - aabb.y;
|
||||
aabb.pMin.x = std::min({p0.x, p1.x, p2.x, p3.x});
|
||||
aabb.pMin.y = std::min({p0.y, p1.y, p2.y, p3.y});
|
||||
aabb.pMax.x = std::max({p0.x, p1.x, p2.x, p3.x});
|
||||
aabb.pMax.y = std::max({p0.y, p1.y, p2.y, p3.y});
|
||||
}
|
||||
|
||||
|
||||
|
|
@ -448,7 +448,7 @@ void WgRenderDataShape::updateMeshes(WgContext& context, const RenderShape &rsha
|
|||
(this->meshGroupStrokesBBox.meshes.count > 0)) {
|
||||
updateAABB(tr);
|
||||
meshDataBBox.bbox(context, pMin, pMax);
|
||||
} else aabb = {0, 0, 0, 0};
|
||||
} else aabb = {{0, 0}, {0, 0}};
|
||||
}
|
||||
|
||||
|
||||
|
|
@ -487,7 +487,7 @@ void WgRenderDataShape::releaseMeshes(WgContext& context)
|
|||
meshGroupShapes.release(context);
|
||||
pMin = {FLT_MAX, FLT_MAX};
|
||||
pMax = {0.0f, 0.0f};
|
||||
aabb = {0, 0, 0, 0};
|
||||
aabb = {{0, 0}, {0, 0}};
|
||||
clips.clear();
|
||||
}
|
||||
|
||||
|
|
@ -600,3 +600,119 @@ void WgRenderDataPicturePool::release(WgContext& context)
|
|||
mPool.clear();
|
||||
mList.clear();
|
||||
}
|
||||
|
||||
//***********************************************************************
|
||||
// WgRenderDataGaussian
|
||||
//***********************************************************************
|
||||
|
||||
void WgRenderDataViewport::update(WgContext& context, const RenderRegion& region) {
|
||||
WgShaderTypeVec4f viewport;
|
||||
viewport.update(region);
|
||||
bool bufferViewportChanged = context.allocateBufferUniform(bufferViewport, &viewport, sizeof(viewport));
|
||||
if (bufferViewportChanged) {
|
||||
context.layouts.releaseBindGroup(bindGroupViewport);
|
||||
bindGroupViewport = context.layouts.createBindGroupBuffer1Un(bufferViewport);
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
void WgRenderDataViewport::release(WgContext& context) {
|
||||
context.releaseBuffer(bufferViewport);
|
||||
context.layouts.releaseBindGroup(bindGroupViewport);
|
||||
}
|
||||
|
||||
//***********************************************************************
|
||||
// WgRenderDataViewportPool
|
||||
//***********************************************************************
|
||||
|
||||
WgRenderDataViewport* WgRenderDataViewportPool::allocate(WgContext& context)
|
||||
{
|
||||
WgRenderDataViewport* renderData{};
|
||||
if (mPool.count > 0) {
|
||||
renderData = mPool.last();
|
||||
mPool.pop();
|
||||
} else {
|
||||
renderData = new WgRenderDataViewport();
|
||||
mList.push(renderData);
|
||||
}
|
||||
return renderData;
|
||||
}
|
||||
|
||||
|
||||
void WgRenderDataViewportPool::free(WgContext& context, WgRenderDataViewport* renderData)
|
||||
{
|
||||
if (renderData) mPool.push(renderData);
|
||||
}
|
||||
|
||||
|
||||
void WgRenderDataViewportPool::release(WgContext& context)
|
||||
{
|
||||
ARRAY_FOREACH(p, mList) {
|
||||
(*p)->release(context);
|
||||
delete(*p);
|
||||
}
|
||||
mPool.clear();
|
||||
mList.clear();
|
||||
}
|
||||
|
||||
//***********************************************************************
|
||||
// WgRenderDataGaussian
|
||||
//***********************************************************************
|
||||
|
||||
void WgRenderDataGaussian::update(WgContext& context, RenderEffectGaussianBlur* gaussian, const Matrix& transform)
|
||||
{
|
||||
assert(gaussian);
|
||||
// compute gaussian blur data
|
||||
WgShaderTypeGaussianBlur gaussianSettings;
|
||||
gaussianSettings.update(gaussian, transform);
|
||||
// update bind group and buffers
|
||||
bool bufferSettingsChanged = context.allocateBufferUniform(bufferSettings, &gaussianSettings.settings, sizeof(gaussianSettings.settings));
|
||||
if (bufferSettingsChanged) {
|
||||
// update bind group
|
||||
context.layouts.releaseBindGroup(bindGroupGaussian);
|
||||
bindGroupGaussian = context.layouts.createBindGroupBuffer1Un(bufferSettings);
|
||||
}
|
||||
level = int(WG_GAUSSIAN_MAX_LEVEL * ((gaussian->quality - 1) * 0.01f)) + 1;
|
||||
extend = gaussianSettings.extend;
|
||||
}
|
||||
|
||||
|
||||
void WgRenderDataGaussian::release(WgContext& context)
|
||||
{
|
||||
context.releaseBuffer(bufferSettings);
|
||||
context.layouts.releaseBindGroup(bindGroupGaussian);
|
||||
}
|
||||
|
||||
//***********************************************************************
|
||||
// WgRenderDataGaussianPool
|
||||
//***********************************************************************
|
||||
|
||||
WgRenderDataGaussian* WgRenderDataGaussianPool::allocate(WgContext& context)
|
||||
{
|
||||
WgRenderDataGaussian* renderData{};
|
||||
if (mPool.count > 0) {
|
||||
renderData = mPool.last();
|
||||
mPool.pop();
|
||||
} else {
|
||||
renderData = new WgRenderDataGaussian();
|
||||
mList.push(renderData);
|
||||
}
|
||||
return renderData;
|
||||
}
|
||||
|
||||
|
||||
void WgRenderDataGaussianPool::free(WgContext& context, WgRenderDataGaussian* renderData)
|
||||
{
|
||||
if (renderData) mPool.push(renderData);
|
||||
}
|
||||
|
||||
|
||||
void WgRenderDataGaussianPool::release(WgContext& context)
|
||||
{
|
||||
ARRAY_FOREACH(p, mList) {
|
||||
(*p)->release(context);
|
||||
delete(*p);
|
||||
}
|
||||
mPool.clear();
|
||||
mList.clear();
|
||||
}
|
||||
|
|
|
|||
|
|
@ -26,6 +26,11 @@
|
|||
#include "tvgWgPipelines.h"
|
||||
#include "tvgWgGeometry.h"
|
||||
|
||||
struct WgAabb {
|
||||
Point pMin{};
|
||||
Point pMax{};
|
||||
};
|
||||
|
||||
struct WgMeshData {
|
||||
WGPUBuffer bufferPosition{};
|
||||
WGPUBuffer bufferTexCoord{};
|
||||
|
|
@ -99,7 +104,7 @@ struct WgRenderDataPaint
|
|||
WGPUBuffer bufferBlendSettings{};
|
||||
WGPUBindGroup bindGroupPaint{};
|
||||
RenderRegion viewport{};
|
||||
RenderRegion aabb{};
|
||||
WgAabb aabb{};
|
||||
float opacity{};
|
||||
Array<WgRenderDataPaint*> clips;
|
||||
|
||||
|
|
@ -167,4 +172,51 @@ public:
|
|||
void release(WgContext& context);
|
||||
};
|
||||
|
||||
struct WgRenderDataViewport
|
||||
{
|
||||
WGPUBindGroup bindGroupViewport{};
|
||||
WGPUBuffer bufferViewport{};
|
||||
|
||||
void update(WgContext& context, const RenderRegion& region);
|
||||
void release(WgContext& context);
|
||||
};
|
||||
|
||||
class WgRenderDataViewportPool {
|
||||
private:
|
||||
// pool contains all created but unused render data for viewport
|
||||
Array<WgRenderDataViewport*> mPool;
|
||||
// list contains all created render data for viewport
|
||||
// to ensure that all created instances will be released
|
||||
Array<WgRenderDataViewport*> mList;
|
||||
public:
|
||||
WgRenderDataViewport* allocate(WgContext& context);
|
||||
void free(WgContext& context, WgRenderDataViewport* renderData);
|
||||
void release(WgContext& context);
|
||||
};
|
||||
|
||||
#define WG_GAUSSIAN_MAX_LEVEL 3
|
||||
struct WgRenderDataGaussian
|
||||
{
|
||||
WGPUBindGroup bindGroupGaussian{};
|
||||
WGPUBuffer bufferSettings{};
|
||||
uint32_t extend{};
|
||||
uint32_t level{};
|
||||
|
||||
void update(WgContext& context, RenderEffectGaussianBlur* gaussian, const Matrix& transform);
|
||||
void release(WgContext& context);
|
||||
};
|
||||
|
||||
class WgRenderDataGaussianPool {
|
||||
private:
|
||||
// pool contains all created but unused render data for gaussian filter
|
||||
Array<WgRenderDataGaussian*> mPool;
|
||||
// list contains all created render data for gaussian filter
|
||||
// to ensure that all created instances will be released
|
||||
Array<WgRenderDataGaussian*> mList;
|
||||
public:
|
||||
WgRenderDataGaussian* allocate(WgContext& context);
|
||||
void free(WgContext& context, WgRenderDataGaussian* renderData);
|
||||
void release(WgContext& context);
|
||||
};
|
||||
|
||||
#endif // _TVG_WG_RENDER_DATA_H_
|
||||
|
|
|
|||
|
|
@ -49,6 +49,8 @@ void WgRenderer::release()
|
|||
// clear render data paint pools
|
||||
mRenderDataShapePool.release(mContext);
|
||||
mRenderDataPicturePool.release(mContext);
|
||||
mRenderDataGaussianPool.release(mContext);
|
||||
mRenderDataViewportPool.release(mContext);
|
||||
WgMeshDataPool::gMeshDataPool->release(mContext);
|
||||
|
||||
// clear render storage pool
|
||||
|
|
@ -182,6 +184,10 @@ bool WgRenderer::postRender()
|
|||
// pop root render storage to the render tree stack
|
||||
mRenderStorageStack.pop();
|
||||
assert(mRenderStorageStack.count == 0);
|
||||
// clear viewport list and store allocated handles to pool
|
||||
ARRAY_FOREACH(p, mRenderDataViewportList)
|
||||
mRenderDataViewportPool.free(mContext, *p);
|
||||
mRenderDataViewportList.clear();
|
||||
return true;
|
||||
}
|
||||
|
||||
|
|
@ -197,7 +203,14 @@ RenderRegion WgRenderer::region(RenderData data)
|
|||
{
|
||||
auto renderData = (WgRenderDataPaint*)data;
|
||||
if (renderData->type() == Type::Shape) {
|
||||
return renderData->aabb;
|
||||
Point v1 = renderData->aabb.pMin;
|
||||
Point v2 = renderData->aabb.pMax;
|
||||
RenderRegion renderRegion;
|
||||
renderRegion.x = static_cast<int32_t>(nearbyint(v1.x));
|
||||
renderRegion.y = static_cast<int32_t>(nearbyint(v1.y));
|
||||
renderRegion.w = static_cast<int32_t>(nearbyint(v2.x)) - renderRegion.x;
|
||||
renderRegion.h = static_cast<int32_t>(nearbyint(v2.y)) - renderRegion.y;
|
||||
return renderRegion;
|
||||
}
|
||||
return { 0, 0, (int32_t)mTargetSurface.w, (int32_t)mTargetSurface.h };
|
||||
}
|
||||
|
|
@ -395,16 +408,24 @@ bool WgRenderer::surfaceConfigure(WGPUSurface surface, WgContext& context, uint3
|
|||
|
||||
RenderCompositor* WgRenderer::target(const RenderRegion& region, TVG_UNUSED ColorSpace cs, TVG_UNUSED CompositionFlag flags)
|
||||
{
|
||||
mCompositorStack.push(new WgCompose);
|
||||
mCompositorStack.last()->aabb = region;
|
||||
return mCompositorStack.last();
|
||||
// create and setup compose data
|
||||
WgCompose* compose = new WgCompose();
|
||||
compose->aabb = region;
|
||||
if (flags & PostProcessing) {
|
||||
compose->aabb = region;
|
||||
compose->rdViewport = mRenderDataViewportPool.allocate(mContext);
|
||||
compose->rdViewport->update(mContext, region);
|
||||
mRenderDataViewportList.push(compose->rdViewport);
|
||||
}
|
||||
mCompositorStack.push(compose);
|
||||
return compose;
|
||||
}
|
||||
|
||||
|
||||
bool WgRenderer::beginComposite(RenderCompositor* cmp, MaskMethod method, uint8_t opacity)
|
||||
{
|
||||
// save current composition settings
|
||||
WgCompose* compose = (WgCompose *)cmp;
|
||||
WgCompose* compose = (WgCompose*)cmp;
|
||||
compose->method = method;
|
||||
compose->opacity = opacity;
|
||||
compose->blend = mBlendMethod;
|
||||
|
|
@ -424,8 +445,8 @@ bool WgRenderer::beginComposite(RenderCompositor* cmp, MaskMethod method, uint8_
|
|||
bool WgRenderer::endComposite(RenderCompositor* cmp)
|
||||
{
|
||||
// get current composition settings
|
||||
WgCompose* comp = (WgCompose *)cmp;
|
||||
// end current render pass
|
||||
WgCompose* comp = (WgCompose*)cmp;
|
||||
// we must to end current render pass to run blend/composition mechanics
|
||||
mCompositor.endRenderPass();
|
||||
// finish scene blending
|
||||
if (comp->method == MaskMethod::None) {
|
||||
|
|
@ -463,30 +484,67 @@ bool WgRenderer::endComposite(RenderCompositor* cmp)
|
|||
}
|
||||
|
||||
|
||||
void WgRenderer::prepare(TVG_UNUSED RenderEffect* effect, TVG_UNUSED const Matrix& transform)
|
||||
void WgRenderer::prepare(RenderEffect* effect, const Matrix& transform)
|
||||
{
|
||||
//TODO: prepare the effect
|
||||
// prepare gaussian blur data
|
||||
if (effect->type == SceneEffect::GaussianBlur) {
|
||||
auto gaussianBlur = (RenderEffectGaussianBlur*)effect;
|
||||
auto renderDataGaussian = (WgRenderDataGaussian*)gaussianBlur->rd;
|
||||
if (!renderDataGaussian) {
|
||||
renderDataGaussian = mRenderDataGaussianPool.allocate(mContext);
|
||||
gaussianBlur->rd = renderDataGaussian;
|
||||
}
|
||||
renderDataGaussian->update(mContext, gaussianBlur, transform);
|
||||
effect->valid = true;
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
bool WgRenderer::region(TVG_UNUSED RenderEffect* effect)
|
||||
bool WgRenderer::region(RenderEffect* effect)
|
||||
{
|
||||
//TODO: Return if the current post effect requires the region expansion
|
||||
if (effect->type == SceneEffect::GaussianBlur) {
|
||||
auto gaussian = (RenderEffectGaussianBlur*)effect;
|
||||
auto renderDataGaussian = (WgRenderDataGaussian*)gaussian->rd;
|
||||
if (gaussian->direction != 2) {
|
||||
gaussian->extend.x = -renderDataGaussian->extend;
|
||||
gaussian->extend.w = +renderDataGaussian->extend * 2;
|
||||
}
|
||||
if (gaussian->direction != 1) {
|
||||
gaussian->extend.y = -renderDataGaussian->extend;
|
||||
gaussian->extend.h = +renderDataGaussian->extend * 2;
|
||||
}
|
||||
return true;
|
||||
}
|
||||
return false;
|
||||
}
|
||||
|
||||
|
||||
bool WgRenderer::render(TVG_UNUSED RenderCompositor* cmp, TVG_UNUSED const RenderEffect* effect, TVG_UNUSED bool direct)
|
||||
bool WgRenderer::render(RenderCompositor* cmp, const RenderEffect* effect, TVG_UNUSED bool direct)
|
||||
{
|
||||
TVGLOG("WG_ENGINE", "SceneEffect(%d) is not supported", (int)effect->type);
|
||||
// we must to end current render pass to resolve ms texture before effect
|
||||
mCompositor.endRenderPass();
|
||||
|
||||
// handle gaussian blur
|
||||
if (effect->type == SceneEffect::GaussianBlur) {
|
||||
WgCompose* comp = (WgCompose*)cmp;
|
||||
WgRenderStorage* dst = mRenderStorageStack.last();
|
||||
RenderEffectGaussianBlur* gaussianBlur = (RenderEffectGaussianBlur*)effect;
|
||||
mCompositor.gaussianBlur(mContext, dst, gaussianBlur, comp);
|
||||
return true;
|
||||
}
|
||||
return false;
|
||||
}
|
||||
|
||||
|
||||
void WgRenderer::dispose(TVG_UNUSED RenderEffect* effect)
|
||||
void WgRenderer::dispose(RenderEffect* effect)
|
||||
{
|
||||
//TODO: dispose the effect
|
||||
}
|
||||
// dispose gaussian blur data
|
||||
if (effect->type == SceneEffect::GaussianBlur) {
|
||||
auto gaussianBlur = (RenderEffectGaussianBlur*)effect;
|
||||
mRenderDataGaussianPool.free(mContext, (WgRenderDataGaussian*)gaussianBlur->rd);
|
||||
gaussianBlur->rd = nullptr;
|
||||
}
|
||||
};
|
||||
|
||||
|
||||
bool WgRenderer::preUpdate()
|
||||
|
|
|
|||
|
|
@ -53,10 +53,10 @@ public:
|
|||
bool beginComposite(RenderCompositor* cmp, MaskMethod method, uint8_t opacity) override;
|
||||
bool endComposite(RenderCompositor* cmp) override;
|
||||
|
||||
void prepare(TVG_UNUSED RenderEffect* effect, TVG_UNUSED const Matrix& transform) override;
|
||||
void prepare(RenderEffect* effect, const Matrix& transform) override;
|
||||
bool region(RenderEffect* effect) override;
|
||||
bool render(RenderCompositor* cmp, const RenderEffect* effect, bool direct) override;
|
||||
void dispose(TVG_UNUSED RenderEffect* effect) override;
|
||||
void dispose(RenderEffect* effect) override;
|
||||
|
||||
static WgRenderer* gen();
|
||||
static bool init(uint32_t threads);
|
||||
|
|
@ -77,6 +77,7 @@ private:
|
|||
WgRenderStorage mRenderStorageRoot;
|
||||
Array<WgCompose*> mCompositorStack;
|
||||
Array<WgRenderStorage*> mRenderStorageStack;
|
||||
Array<WgRenderDataViewport*> mRenderDataViewportList;
|
||||
|
||||
// render storage pool
|
||||
WgRenderStoragePool mRenderStoragePool;
|
||||
|
|
@ -84,6 +85,8 @@ private:
|
|||
// render data paint pools
|
||||
WgRenderDataShapePool mRenderDataShapePool;
|
||||
WgRenderDataPicturePool mRenderDataPicturePool;
|
||||
WgRenderDataGaussianPool mRenderDataGaussianPool;
|
||||
WgRenderDataViewportPool mRenderDataViewportPool;
|
||||
|
||||
// rendering context
|
||||
WgContext mContext;
|
||||
|
|
|
|||
|
|
@ -719,4 +719,96 @@ fn cs_main(@builtin(global_invocation_id) id: vec3u) {
|
|||
let colorMsk1 = textureLoad(imageMsk1, id.xy);
|
||||
textureStore(imageTrg, id.xy, colorMsk0 * colorMsk1);
|
||||
}
|
||||
)";
|
||||
)";
|
||||
|
||||
const char* cShaderSrc_GaussianBlur_Horz = R"(
|
||||
@group(0) @binding(0) var imageSrc : texture_storage_2d<rgba8unorm, read>;
|
||||
@group(1) @binding(0) var imageDst : texture_storage_2d<rgba8unorm, write>;
|
||||
@group(2) @binding(0) var<uniform> settings: vec4f;
|
||||
@group(3) @binding(0) var<uniform> viewport: vec4f;
|
||||
|
||||
fn gaussian(x: f32, sigma: f32) -> f32 {
|
||||
let a = 0.39894f / sigma;
|
||||
let b = -(x * x) / (2.0 * sigma * sigma);
|
||||
return a * exp(b);
|
||||
}
|
||||
|
||||
@compute @workgroup_size(16, 16)
|
||||
fn cs_main(@builtin(global_invocation_id) id: vec3u) {
|
||||
// id conversion
|
||||
let iid = vec2i(id.xy);
|
||||
// viewport decode
|
||||
let xmin = i32(viewport.x);
|
||||
let ymin = i32(viewport.y);
|
||||
let xmax = i32(viewport.z);
|
||||
let ymax = i32(viewport.w);
|
||||
// settings decode
|
||||
let sigma = settings.x;
|
||||
let scale = settings.y;
|
||||
let size = i32(settings.z);
|
||||
|
||||
// draw borders points outside of viewport
|
||||
if ((iid.x < xmin) || (iid.x > xmax) || (iid.y < ymin) || (iid.y > ymax)) { return; }
|
||||
|
||||
// apply filter
|
||||
var weight = gaussian(0.0, sigma);
|
||||
var color = weight * textureLoad(imageSrc, id.xy);
|
||||
var sum = weight;
|
||||
for (var i: i32 = 1; i < size; i++) {
|
||||
let ii = i32(f32(i) * scale);
|
||||
let idneg = vec2i(clamp(iid.x - ii, xmin, xmax), iid.y);
|
||||
let idpos = vec2i(clamp(iid.x + ii, xmin, xmax), iid.y);
|
||||
weight = gaussian(f32(i) * scale, sigma);
|
||||
color += (weight * textureLoad(imageSrc, vec2u(idneg)));
|
||||
color += (weight * textureLoad(imageSrc, vec2u(idpos)));
|
||||
sum += (2.0 * weight);
|
||||
}
|
||||
textureStore(imageDst, id.xy, color / sum);
|
||||
}
|
||||
)";
|
||||
|
||||
const char* cShaderSrc_GaussianBlur_Vert = R"(
|
||||
@group(0) @binding(0) var imageSrc : texture_storage_2d<rgba8unorm, read>;
|
||||
@group(1) @binding(0) var imageDst : texture_storage_2d<rgba8unorm, write>;
|
||||
@group(2) @binding(0) var<uniform> settings: vec4f;
|
||||
@group(3) @binding(0) var<uniform> viewport: vec4f;
|
||||
|
||||
fn gaussian(x: f32, sigma: f32) -> f32 {
|
||||
let a = 0.39894f / sigma;
|
||||
let b = -(x * x) / (2.0 * sigma * sigma);
|
||||
return a * exp(b);
|
||||
}
|
||||
|
||||
@compute @workgroup_size(16, 16)
|
||||
fn cs_main(@builtin(global_invocation_id) id: vec3u) {
|
||||
// id conversion
|
||||
let iid = vec2i(id.xy);
|
||||
// viewport decode
|
||||
let xmin = i32(viewport.x);
|
||||
let ymin = i32(viewport.y);
|
||||
let xmax = i32(viewport.z);
|
||||
let ymax = i32(viewport.w);
|
||||
// settings decode
|
||||
let sigma = settings.x;
|
||||
let scale = settings.y;
|
||||
let size = i32(settings.z);
|
||||
|
||||
// draw borders points outside of viewport
|
||||
if ((iid.x < xmin) || (iid.x > xmax) || (iid.y < ymin) || (iid.y > ymax)) { return; }
|
||||
|
||||
// apply filter
|
||||
var weight = gaussian(0.0, sigma);
|
||||
var color = weight * textureLoad(imageSrc, id.xy);
|
||||
var sum = weight;
|
||||
for (var i: i32 = 1; i < size; i++) {
|
||||
let ii = i32(f32(i) * scale);
|
||||
let idneg = vec2i(iid.x, clamp(iid.y - ii, ymin, ymax));
|
||||
let idpos = vec2i(iid.x, clamp(iid.y + ii, ymin, ymax));
|
||||
weight = gaussian(f32(i) * scale, sigma);
|
||||
color += (weight * textureLoad(imageSrc, vec2u(idneg)));
|
||||
color += (weight * textureLoad(imageSrc, vec2u(idpos)));
|
||||
sum += (2.0 * weight);
|
||||
}
|
||||
textureStore(imageDst, id.xy, color / sum);
|
||||
}
|
||||
)";
|
||||
|
|
|
|||
|
|
@ -44,7 +44,9 @@ extern const char* cShaderSrc_Scene_Compose;
|
|||
// shaders blit
|
||||
extern const char* cShaderSrc_Blit;
|
||||
|
||||
// compute shader sources: blend, compose and merge path
|
||||
// compute shader sources: effects
|
||||
extern const char* cShaderSrc_MergeMasks;
|
||||
extern const char* cShaderSrc_GaussianBlur_Vert;
|
||||
extern const char* cShaderSrc_GaussianBlur_Horz;
|
||||
|
||||
#endif // _TVG_WG_SHEDER_SRC_H_
|
||||
|
|
|
|||
|
|
@ -116,6 +116,14 @@ void WgShaderTypeVec4f::update(const RenderColor& c)
|
|||
vec[3] = c.a / 255.0f; // alpha
|
||||
}
|
||||
|
||||
void WgShaderTypeVec4f::update(const RenderRegion& r)
|
||||
{
|
||||
vec[0] = r.x; // left
|
||||
vec[1] = r.y; // top
|
||||
vec[2] = r.x + r.w - 1; // right
|
||||
vec[3] = r.y + r.h - 1; // bottom
|
||||
}
|
||||
|
||||
//************************************************************************
|
||||
// WgShaderTypeGradient
|
||||
//************************************************************************
|
||||
|
|
@ -187,3 +195,18 @@ void WgShaderTypeGradient::updateTexData(const Fill::ColorStop* stops, uint32_t
|
|||
texData[ti * 4 + 3] = colorStopLast.a;
|
||||
}
|
||||
}
|
||||
|
||||
//************************************************************************
|
||||
// WgShaderTypeGaussianBlur
|
||||
//************************************************************************
|
||||
|
||||
void WgShaderTypeGaussianBlur::update(const RenderEffectGaussianBlur* gaussian, const Matrix& transform)
|
||||
{
|
||||
assert(gaussian);
|
||||
const float sigma = gaussian->sigma;
|
||||
const float scale = std::sqrt(transform.e11 * transform.e11 + transform.e12 * transform.e12);
|
||||
settings[0] = sigma;
|
||||
settings[1] = scale;
|
||||
settings[2] = 2 * sigma * scale; // kernel size
|
||||
extend = settings[2] * 2;
|
||||
}
|
||||
|
|
@ -52,6 +52,7 @@ struct WgShaderTypeVec4f
|
|||
WgShaderTypeVec4f(const RenderColor& c);
|
||||
void update(const ColorSpace colorSpace, uint8_t o);
|
||||
void update(const RenderColor& c);
|
||||
void update(const RenderRegion& r);
|
||||
};
|
||||
|
||||
// sampler, texture, vec4f
|
||||
|
|
@ -66,4 +67,13 @@ struct WgShaderTypeGradient
|
|||
void updateTexData(const Fill::ColorStop* stops, uint32_t stopCnt);
|
||||
};
|
||||
|
||||
// gaussian settings: sigma, scale, extend
|
||||
struct WgShaderTypeGaussianBlur
|
||||
{
|
||||
float settings[4]{}; // [0]: sigma, [1]: scale, [2]: kernel size, [3]: unused
|
||||
uint32_t extend{};
|
||||
|
||||
void update(const RenderEffectGaussianBlur* gaussian, const Matrix& transform);
|
||||
};
|
||||
|
||||
#endif // _TVG_WG_SHADER_TYPES_H_
|
||||
|
|
|
|||
Loading…
Add table
Reference in a new issue