wg_engine: ClipPath support

[issues 1479: ClipPath](#1479)

Supports ClipPath composition.
Clip path composition is an only composition type who doesn't ignore blend method.
Clip path is a combination of composition approach and blend approach using compute shader
This commit is contained in:
Sergii Liebodkin 2024-07-13 08:27:17 +00:00 committed by Hermet Park
parent f06127ab93
commit 2c948a33d3
12 changed files with 530 additions and 45 deletions

View file

@ -36,11 +36,15 @@ WGPUBindGroupLayout WgBindGroupTexture::layout = nullptr;
WGPUBindGroupLayout WgBindGroupTextureStorageRgba::layout = nullptr; WGPUBindGroupLayout WgBindGroupTextureStorageRgba::layout = nullptr;
WGPUBindGroupLayout WgBindGroupTextureStorageBgra::layout = nullptr; WGPUBindGroupLayout WgBindGroupTextureStorageBgra::layout = nullptr;
WGPUBindGroupLayout WgBindGroupTextureSampled::layout = nullptr; WGPUBindGroupLayout WgBindGroupTextureSampled::layout = nullptr;
WGPUBindGroupLayout WgBindGroupTexMaskCompose::layout = nullptr;
WGPUBindGroupLayout WgBindGroupTexComposeBlend::layout = nullptr; WGPUBindGroupLayout WgBindGroupTexComposeBlend::layout = nullptr;
WGPUBindGroupLayout WgBindGroupOpacity::layout = nullptr; WGPUBindGroupLayout WgBindGroupOpacity::layout = nullptr;
WGPUBindGroupLayout WgBindGroupBlendMethod::layout = nullptr; WGPUBindGroupLayout WgBindGroupBlendMethod::layout = nullptr;
WGPUBindGroupLayout WgBindGroupCompositeMethod::layout = nullptr; WGPUBindGroupLayout WgBindGroupCompositeMethod::layout = nullptr;
///////////////////////////////////////////////////////////////////////////////
// WgBindGroupCanvas
///////////////////////////////////////////////////////////////////////////////
WGPUBindGroupLayout WgBindGroupCanvas::getLayout(WGPUDevice device) WGPUBindGroupLayout WgBindGroupCanvas::getLayout(WGPUDevice device)
{ {
@ -78,6 +82,9 @@ void WgBindGroupCanvas::release()
releaseBuffer(uBufferViewMat); releaseBuffer(uBufferViewMat);
} }
///////////////////////////////////////////////////////////////////////////////
// WgBindGroupPaint
///////////////////////////////////////////////////////////////////////////////
WGPUBindGroupLayout WgBindGroupPaint::getLayout(WGPUDevice device) WGPUBindGroupLayout WgBindGroupPaint::getLayout(WGPUDevice device)
{ {
@ -123,6 +130,9 @@ void WgBindGroupPaint::release()
releaseBuffer(uBufferModelMat); releaseBuffer(uBufferModelMat);
} }
///////////////////////////////////////////////////////////////////////////////
// WgBindGroupSolidColor
///////////////////////////////////////////////////////////////////////////////
WGPUBindGroupLayout WgBindGroupSolidColor::getLayout(WGPUDevice device) WGPUBindGroupLayout WgBindGroupSolidColor::getLayout(WGPUDevice device)
{ {
@ -163,6 +173,9 @@ void WgBindGroupSolidColor::release()
releaseBuffer(uBufferSolidColor); releaseBuffer(uBufferSolidColor);
} }
///////////////////////////////////////////////////////////////////////////////
// WgBindGroupLinearGradient
///////////////////////////////////////////////////////////////////////////////
WGPUBindGroupLayout WgBindGroupLinearGradient::getLayout(WGPUDevice device) WGPUBindGroupLayout WgBindGroupLinearGradient::getLayout(WGPUDevice device)
{ {
@ -203,6 +216,9 @@ void WgBindGroupLinearGradient::release()
releaseBuffer(uBufferLinearGradient); releaseBuffer(uBufferLinearGradient);
} }
///////////////////////////////////////////////////////////////////////////////
// WgBindGroupRadialGradient
///////////////////////////////////////////////////////////////////////////////
WGPUBindGroupLayout WgBindGroupRadialGradient::getLayout(WGPUDevice device) WGPUBindGroupLayout WgBindGroupRadialGradient::getLayout(WGPUDevice device)
{ {
@ -243,6 +259,9 @@ void WgBindGroupRadialGradient::release()
releaseBindGroup(mBindGroup); releaseBindGroup(mBindGroup);
} }
///////////////////////////////////////////////////////////////////////////////
// WgBindGroupPicture
///////////////////////////////////////////////////////////////////////////////
WGPUBindGroupLayout WgBindGroupPicture::getLayout(WGPUDevice device) WGPUBindGroupLayout WgBindGroupPicture::getLayout(WGPUDevice device)
{ {
@ -280,6 +299,9 @@ void WgBindGroupPicture::release()
releaseBindGroup(mBindGroup); releaseBindGroup(mBindGroup);
} }
///////////////////////////////////////////////////////////////////////////////
// WgBindGroupTexture
///////////////////////////////////////////////////////////////////////////////
WGPUBindGroupLayout WgBindGroupTexture::getLayout(WGPUDevice device) WGPUBindGroupLayout WgBindGroupTexture::getLayout(WGPUDevice device)
{ {
@ -315,6 +337,9 @@ void WgBindGroupTexture::release()
releaseBindGroup(mBindGroup); releaseBindGroup(mBindGroup);
} }
///////////////////////////////////////////////////////////////////////////////
// WgBindGroupTextureStorageRgba
///////////////////////////////////////////////////////////////////////////////
WGPUBindGroupLayout WgBindGroupTextureStorageRgba::getLayout(WGPUDevice device) WGPUBindGroupLayout WgBindGroupTextureStorageRgba::getLayout(WGPUDevice device)
{ {
@ -350,6 +375,9 @@ void WgBindGroupTextureStorageRgba::release()
releaseBindGroup(mBindGroup); releaseBindGroup(mBindGroup);
} }
///////////////////////////////////////////////////////////////////////////////
// WgBindGroupTextureStorageBgra
///////////////////////////////////////////////////////////////////////////////
WGPUBindGroupLayout WgBindGroupTextureStorageBgra::getLayout(WGPUDevice device) WGPUBindGroupLayout WgBindGroupTextureStorageBgra::getLayout(WGPUDevice device)
{ {
@ -385,6 +413,9 @@ void WgBindGroupTextureStorageBgra::release()
releaseBindGroup(mBindGroup); releaseBindGroup(mBindGroup);
} }
///////////////////////////////////////////////////////////////////////////////
// WgBindGroupTextureSampled
///////////////////////////////////////////////////////////////////////////////
WGPUBindGroupLayout WgBindGroupTextureSampled::getLayout(WGPUDevice device) WGPUBindGroupLayout WgBindGroupTextureSampled::getLayout(WGPUDevice device)
{ {
@ -422,6 +453,9 @@ void WgBindGroupTextureSampled::release()
releaseBindGroup(mBindGroup); releaseBindGroup(mBindGroup);
} }
///////////////////////////////////////////////////////////////////////////////
// WgBindGroupTexComposeBlend
///////////////////////////////////////////////////////////////////////////////
WGPUBindGroupLayout WgBindGroupTexComposeBlend::getLayout(WGPUDevice device) WGPUBindGroupLayout WgBindGroupTexComposeBlend::getLayout(WGPUDevice device)
{ {
@ -461,7 +495,49 @@ void WgBindGroupTexComposeBlend::release()
releaseBindGroup(mBindGroup); releaseBindGroup(mBindGroup);
} }
///////////////////////////////////////////////////////////////////////////////
// WgBindGroupTexMaskCompose
///////////////////////////////////////////////////////////////////////////////
WGPUBindGroupLayout WgBindGroupTexMaskCompose::getLayout(WGPUDevice device)
{
if (layout) return layout;
const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] {
makeBindGroupLayoutEntryStorage(0, WGPUStorageTextureAccess_ReadOnly, WGPUTextureFormat_RGBA8Unorm),
makeBindGroupLayoutEntryStorage(1, WGPUStorageTextureAccess_ReadWrite, WGPUTextureFormat_RGBA8Unorm)
};
layout = createBindGroupLayout(device, bindGroupLayoutEntries, 2);
assert(layout);
return layout;
}
void WgBindGroupTexMaskCompose::releaseLayout()
{
releaseBindGroupLayout(layout);
}
void WgBindGroupTexMaskCompose::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexMsk0, WGPUTextureView uTexMsk1)
{
release();
const WGPUBindGroupEntry bindGroupEntries[] {
makeBindGroupEntryTextureView(0, uTexMsk0),
makeBindGroupEntryTextureView(1, uTexMsk1),
};
mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 2);
assert(mBindGroup);
}
void WgBindGroupTexMaskCompose::release()
{
releaseBindGroup(mBindGroup);
}
///////////////////////////////////////////////////////////////////////////////
// WgBindGroupOpacity
///////////////////////////////////////////////////////////////////////////////
WGPUBindGroupLayout WgBindGroupOpacity::getLayout(WGPUDevice device) WGPUBindGroupLayout WgBindGroupOpacity::getLayout(WGPUDevice device)
{ {
@ -500,6 +576,9 @@ void WgBindGroupOpacity::release()
releaseBindGroup(mBindGroup); releaseBindGroup(mBindGroup);
} }
///////////////////////////////////////////////////////////////////////////////
// WgBindGroupBlendMethod
///////////////////////////////////////////////////////////////////////////////
WGPUBindGroupLayout WgBindGroupBlendMethod::getLayout(WGPUDevice device) WGPUBindGroupLayout WgBindGroupBlendMethod::getLayout(WGPUDevice device)
{ {
@ -538,6 +617,9 @@ void WgBindGroupBlendMethod::release()
releaseBindGroup(mBindGroup); releaseBindGroup(mBindGroup);
} }
///////////////////////////////////////////////////////////////////////////////
// WgBindGroupCompositeMethod
///////////////////////////////////////////////////////////////////////////////
WGPUBindGroupLayout WgBindGroupCompositeMethod::getLayout(WGPUDevice device) WGPUBindGroupLayout WgBindGroupCompositeMethod::getLayout(WGPUDevice device)
{ {

View file

@ -170,6 +170,20 @@ struct WgBindGroupTexComposeBlend : public WgBindGroup
}; };
// @group(0)
struct WgBindGroupTexMaskCompose : public WgBindGroup
{
static WGPUBindGroupLayout layout;
static WGPUBindGroupLayout getLayout(WGPUDevice device);
static void releaseLayout();
void initialize(WGPUDevice device, WGPUQueue queue,
WGPUTextureView uTexMsk0,
WGPUTextureView uTexMsk1);
void release();
};
// @group(1 or 2) // @group(1 or 2)
struct WgBindGroupOpacity : public WgBindGroup struct WgBindGroupOpacity : public WgBindGroup
{ {

View file

@ -49,7 +49,7 @@ void WgPipelineFillShapeWinding::initialize(WGPUDevice device)
WGPUCompareFunction stencilFunctionBack = WGPUCompareFunction_Always; WGPUCompareFunction stencilFunctionBack = WGPUCompareFunction_Always;
WGPUStencilOperation stencilOperationBack = WGPUStencilOperation_DecrementWrap; WGPUStencilOperation stencilOperationBack = WGPUStencilOperation_DecrementWrap;
// sheder source and labels // shader source and labels
auto shaderSource = cShaderSource_PipelineFill; auto shaderSource = cShaderSource_PipelineFill;
auto shaderLabel = "The shader fill"; auto shaderLabel = "The shader fill";
auto pipelineLabel = "The render pipeline fill shape winding"; auto pipelineLabel = "The render pipeline fill shape winding";
@ -83,7 +83,7 @@ void WgPipelineFillShapeEvenOdd::initialize(WGPUDevice device)
WGPUCompareFunction stencilFunctionBack = WGPUCompareFunction_Always; WGPUCompareFunction stencilFunctionBack = WGPUCompareFunction_Always;
WGPUStencilOperation stencilOperationBack = WGPUStencilOperation_Invert; WGPUStencilOperation stencilOperationBack = WGPUStencilOperation_Invert;
// sheder source and labels // shader source and labels
auto shaderSource = cShaderSource_PipelineFill; auto shaderSource = cShaderSource_PipelineFill;
auto shaderLabel = "The shader fill"; auto shaderLabel = "The shader fill";
auto pipelineLabel = "The render pipeline fill shape Even Odd"; auto pipelineLabel = "The render pipeline fill shape Even Odd";
@ -115,7 +115,7 @@ void WgPipelineFillStroke::initialize(WGPUDevice device)
WGPUCompareFunction stencilFunction = WGPUCompareFunction_Always; WGPUCompareFunction stencilFunction = WGPUCompareFunction_Always;
WGPUStencilOperation stencilOperation = WGPUStencilOperation_Replace; WGPUStencilOperation stencilOperation = WGPUStencilOperation_Replace;
// sheder source and labels // shader source and labels
auto shaderSource = cShaderSource_PipelineFill; auto shaderSource = cShaderSource_PipelineFill;
auto shaderLabel = "The shader fill"; auto shaderLabel = "The shader fill";
auto pipelineLabel = "The render pipeline fill stroke"; auto pipelineLabel = "The render pipeline fill stroke";
@ -129,6 +129,38 @@ void WgPipelineFillStroke::initialize(WGPUDevice device)
} }
void WgPipelineClipMask::initialize(WGPUDevice device)
{
// vertex and buffers settings
WGPUVertexAttribute vertexAttributesPos = { WGPUVertexFormat_Float32x2, sizeof(float) * 0, 0 };
WGPUVertexBufferLayout vertexBufferLayouts[] = {
makeVertexBufferLayout(&vertexAttributesPos, 1, sizeof(float) * 2)
};
// bind groups
WGPUBindGroupLayout bindGroupLayouts[] = {
WgBindGroupCanvas::getLayout(device),
WgBindGroupPaint::getLayout(device)
};
// stencil function
WGPUCompareFunction stencilFunction = WGPUCompareFunction_NotEqual;
WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero;
// shader source and labels
auto shaderSource = cShaderSource_PipelineFill;
auto shaderLabel = "The shader fill";
auto pipelineLabel = "The render pipeline clip mask";
// allocate all pipeline handles
allocate(device, WgPipelineBlendType::SrcOver, WGPUColorWriteMask_All,
vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts),
bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts),
stencilFunction, stencilOperation, stencilFunction, stencilOperation,
shaderSource, shaderLabel, pipelineLabel);
}
void WgPipelineSolid::initialize(WGPUDevice device, WgPipelineBlendType blendType) void WgPipelineSolid::initialize(WGPUDevice device, WgPipelineBlendType blendType)
{ {
// vertex and buffers settings // vertex and buffers settings
@ -148,7 +180,7 @@ void WgPipelineSolid::initialize(WGPUDevice device, WgPipelineBlendType blendTyp
WGPUCompareFunction stencilFunction = WGPUCompareFunction_NotEqual; WGPUCompareFunction stencilFunction = WGPUCompareFunction_NotEqual;
WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero; WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero;
// sheder source and labels // shader source and labels
auto shaderSource = cShaderSource_PipelineSolid; auto shaderSource = cShaderSource_PipelineSolid;
auto shaderLabel = "The shader solid color"; auto shaderLabel = "The shader solid color";
auto pipelineLabel = "The render pipeline solid color"; auto pipelineLabel = "The render pipeline solid color";
@ -181,7 +213,7 @@ void WgPipelineLinear::initialize(WGPUDevice device, WgPipelineBlendType blendTy
WGPUCompareFunction stencilFunction = WGPUCompareFunction_NotEqual; WGPUCompareFunction stencilFunction = WGPUCompareFunction_NotEqual;
WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero; WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero;
// sheder source and labels // shader source and labels
auto shaderSource = cShaderSource_PipelineLinear; auto shaderSource = cShaderSource_PipelineLinear;
auto shaderLabel = "The shader linear gradient"; auto shaderLabel = "The shader linear gradient";
auto pipelineLabel = "The render pipeline linear gradient"; auto pipelineLabel = "The render pipeline linear gradient";
@ -214,7 +246,7 @@ void WgPipelineRadial::initialize(WGPUDevice device, WgPipelineBlendType blendTy
WGPUCompareFunction stencilFunction = WGPUCompareFunction_NotEqual; WGPUCompareFunction stencilFunction = WGPUCompareFunction_NotEqual;
WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero; WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero;
// sheder source and labels // shader source and labels
auto shaderSource = cShaderSource_PipelineRadial; auto shaderSource = cShaderSource_PipelineRadial;
auto shaderLabel = "The shader radial gradient"; auto shaderLabel = "The shader radial gradient";
auto pipelineLabel = "The render pipeline radial gradient"; auto pipelineLabel = "The render pipeline radial gradient";
@ -249,7 +281,7 @@ void WgPipelineImage::initialize(WGPUDevice device, WgPipelineBlendType blendTyp
WGPUCompareFunction stencilFunction = WGPUCompareFunction_Always; WGPUCompareFunction stencilFunction = WGPUCompareFunction_Always;
WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero; WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero;
// sheder source and labels // shader source and labels
auto shaderSource = cShaderSource_PipelineImage; auto shaderSource = cShaderSource_PipelineImage;
auto shaderLabel = "The shader image"; auto shaderLabel = "The shader image";
auto pipelineLabel = "The render pipeline image"; auto pipelineLabel = "The render pipeline image";
@ -273,7 +305,7 @@ void WgPipelineClear::initialize(WGPUDevice device)
WgBindGroupTextureStorageRgba::getLayout(device) WgBindGroupTextureStorageRgba::getLayout(device)
}; };
// sheder source and labels // shader source and labels
auto shaderSource = cShaderSource_PipelineComputeClear; auto shaderSource = cShaderSource_PipelineComputeClear;
auto shaderLabel = "The compute shader clear"; auto shaderLabel = "The compute shader clear";
auto pipelineLabel = "The compute pipeline clear"; auto pipelineLabel = "The compute pipeline clear";
@ -295,7 +327,7 @@ void WgPipelineBlend::initialize(WGPUDevice device, const char *shaderSource)
WgBindGroupOpacity::getLayout(device) WgBindGroupOpacity::getLayout(device)
}; };
// sheder source and labels // shader source and labels
auto shaderLabel = "The compute shader blend"; auto shaderLabel = "The compute shader blend";
auto pipelineLabel = "The compute pipeline blend"; auto pipelineLabel = "The compute pipeline blend";
@ -306,7 +338,46 @@ void WgPipelineBlend::initialize(WGPUDevice device, const char *shaderSource)
} }
void WgPipelineComposeBlend::initialize(WGPUDevice device) void WgPipelineBlendMask::initialize(WGPUDevice device, const char *shaderSource)
{
// bind groups and layouts
WGPUBindGroupLayout bindGroupLayouts[] = {
WgBindGroupTexComposeBlend::getLayout(device),
WgBindGroupBlendMethod::getLayout(device),
WgBindGroupOpacity::getLayout(device)
};
// shader source and labels
auto shaderLabel = "The compute shader blend mask";
auto pipelineLabel = "The compute pipeline blend mask";
// allocate all pipeline handles
allocate(device,
bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts),
shaderSource, shaderLabel, pipelineLabel);
}
void WgPipelineMaskCompose::initialize(WGPUDevice device)
{
// bind groups and layouts
WGPUBindGroupLayout bindGroupLayouts[] = {
WgBindGroupTexMaskCompose::getLayout(device),
};
// shader source and labels
auto shaderSource = cShaderSource_PipelineComputeMaskCompose;
auto shaderLabel = "The compute shader mask compose";
auto pipelineLabel = "The compute pipeline mask compose";
// allocate all pipeline handles
allocate(device,
bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts),
shaderSource, shaderLabel, pipelineLabel);
}
void WgPipelineCompose::initialize(WGPUDevice device)
{ {
// bind groups and layouts // bind groups and layouts
WGPUBindGroupLayout bindGroupLayouts[] = { WGPUBindGroupLayout bindGroupLayouts[] = {
@ -316,8 +387,8 @@ void WgPipelineComposeBlend::initialize(WGPUDevice device)
WgBindGroupOpacity::getLayout(device) WgBindGroupOpacity::getLayout(device)
}; };
// sheder source and labels // shader source and labels
auto shaderSource = cShaderSource_PipelineComputeComposeBlend; auto shaderSource = cShaderSource_PipelineComputeCompose;
auto shaderLabel = "The compute shader compose blend"; auto shaderLabel = "The compute shader compose blend";
auto pipelineLabel = "The compute pipeline compose blend"; auto pipelineLabel = "The compute pipeline compose blend";
@ -336,7 +407,7 @@ void WgPipelineAntiAliasing::initialize(WGPUDevice device)
WgBindGroupTextureStorageBgra::getLayout(device) WgBindGroupTextureStorageBgra::getLayout(device)
}; };
// sheder source and labels // shader source and labels
auto shaderSource = cShaderSource_PipelineComputeAntiAlias; auto shaderSource = cShaderSource_PipelineComputeAntiAlias;
auto shaderLabel = "The compute shader anti-aliasing"; auto shaderLabel = "The compute shader anti-aliasing";
auto pipelineLabel = "The compute pipeline anti-aliasing"; auto pipelineLabel = "The compute pipeline anti-aliasing";
@ -357,6 +428,7 @@ void WgPipelines::initialize(WgContext& context)
fillShapeWinding.initialize(context.device); fillShapeWinding.initialize(context.device);
fillShapeEvenOdd.initialize(context.device); fillShapeEvenOdd.initialize(context.device);
fillStroke.initialize(context.device); fillStroke.initialize(context.device);
clipMask.initialize(context.device);
for (uint8_t type = (uint8_t)WgPipelineBlendType::SrcOver; type <= (uint8_t)WgPipelineBlendType::Custom; type++) { for (uint8_t type = (uint8_t)WgPipelineBlendType::SrcOver; type <= (uint8_t)WgPipelineBlendType::Custom; type++) {
solid[type].initialize(context.device, (WgPipelineBlendType)type); solid[type].initialize(context.device, (WgPipelineBlendType)type);
linear[type].initialize(context.device, (WgPipelineBlendType)type); linear[type].initialize(context.device, (WgPipelineBlendType)type);
@ -368,7 +440,11 @@ void WgPipelines::initialize(WgContext& context)
computeBlendSolid.initialize(context.device, cShaderSource_PipelineComputeBlendSolid); computeBlendSolid.initialize(context.device, cShaderSource_PipelineComputeBlendSolid);
computeBlendGradient.initialize(context.device, cShaderSource_PipelineComputeBlendGradient); computeBlendGradient.initialize(context.device, cShaderSource_PipelineComputeBlendGradient);
computeBlendImage.initialize(context.device, cShaderSource_PipelineComputeBlendImage); computeBlendImage.initialize(context.device, cShaderSource_PipelineComputeBlendImage);
computeComposeBlend.initialize(context.device); computeBlendSolidMask.initialize(context.device, cShaderSource_PipelineComputeBlendSolidMask);
computeBlendGradientMask.initialize(context.device, cShaderSource_PipelineComputeBlendGradientMask);
computeBlendImageMask.initialize(context.device, cShaderSource_PipelineComputeBlendImageMask);
computeMaskCompose.initialize(context.device);
computeCompose.initialize(context.device);
computeAntiAliasing.initialize(context.device); computeAntiAliasing.initialize(context.device);
// store pipelines to context // store pipelines to context
context.pipelines = this; context.pipelines = this;
@ -377,7 +453,8 @@ void WgPipelines::initialize(WgContext& context)
void WgPipelines::release() void WgPipelines::release()
{ {
WgBindGroupTexComposeBlend::layout = nullptr; WgBindGroupTexMaskCompose::releaseLayout();
WgBindGroupTexComposeBlend::releaseLayout();
WgBindGroupTextureSampled::releaseLayout(); WgBindGroupTextureSampled::releaseLayout();
WgBindGroupTextureStorageBgra::releaseLayout(); WgBindGroupTextureStorageBgra::releaseLayout();
WgBindGroupTextureStorageRgba::releaseLayout(); WgBindGroupTextureStorageRgba::releaseLayout();
@ -391,7 +468,11 @@ void WgPipelines::release()
WgBindGroupCanvas::releaseLayout(); WgBindGroupCanvas::releaseLayout();
// compute pipelines // compute pipelines
computeAntiAliasing.release(); computeAntiAliasing.release();
computeComposeBlend.release(); computeCompose.release();
computeMaskCompose.release();
computeBlendImageMask.release();
computeBlendGradientMask.release();
computeBlendSolidMask.release();
computeBlendImage.release(); computeBlendImage.release();
computeBlendGradient.release(); computeBlendGradient.release();
computeBlendSolid.release(); computeBlendSolid.release();
@ -403,6 +484,7 @@ void WgPipelines::release()
linear[type].release(); linear[type].release();
solid[type].release(); solid[type].release();
} }
clipMask.release();
fillStroke.release(); fillStroke.release();
fillShapeEvenOdd.release(); fillShapeEvenOdd.release();
fillShapeWinding.release(); fillShapeWinding.release();

View file

@ -62,6 +62,17 @@ struct WgPipelineFillStroke: public WgRenderPipeline
} }
}; };
struct WgPipelineClipMask: public WgRenderPipeline
{
void initialize(WGPUDevice device) override;
void use(WGPURenderPassEncoder encoder, WgBindGroupCanvas& groupCanvas,WgBindGroupPaint& groupPaint)
{
set(encoder);
groupCanvas.set(encoder, 0);
groupPaint.set(encoder, 1);
}
};
struct WgPipelineSolid: public WgRenderPipeline struct WgPipelineSolid: public WgRenderPipeline
{ {
void initialize(WGPUDevice device) override {} void initialize(WGPUDevice device) override {}
@ -144,7 +155,32 @@ struct WgPipelineBlend: public WgComputePipeline
}; };
struct WgPipelineComposeBlend: public WgComputePipeline struct WgPipelineBlendMask: public WgComputePipeline
{
void initialize(WGPUDevice device) override { assert(false); };
void initialize(WGPUDevice device, const char *shaderSource);
void use(WGPUComputePassEncoder encoder, WgBindGroupTexComposeBlend& groupTexs, WgBindGroupBlendMethod& blendMethod, WgBindGroupOpacity& groupOpacity)
{
set(encoder);
groupTexs.set(encoder, 0);
blendMethod.set(encoder, 1);
groupOpacity.set(encoder, 2);
}
};
struct WgPipelineMaskCompose: public WgComputePipeline
{
void initialize(WGPUDevice device) override;
void use(WGPUComputePassEncoder encoder, WgBindGroupTexMaskCompose& groupTexs)
{
set(encoder);
groupTexs.set(encoder, 0);
}
};
struct WgPipelineCompose: public WgComputePipeline
{ {
void initialize(WGPUDevice device) override; void initialize(WGPUDevice device) override;
void use(WGPUComputePassEncoder encoder, WgBindGroupTexComposeBlend& groupTexs, WgBindGroupCompositeMethod& groupComposeMethod, WgBindGroupBlendMethod& groupBlendMethod, WgBindGroupOpacity& groupOpacity) void use(WGPUComputePassEncoder encoder, WgBindGroupTexComposeBlend& groupTexs, WgBindGroupCompositeMethod& groupComposeMethod, WgBindGroupBlendMethod& groupBlendMethod, WgBindGroupOpacity& groupOpacity)
@ -180,6 +216,7 @@ struct WgPipelines
WgPipelineFillShapeEvenOdd fillShapeEvenOdd; WgPipelineFillShapeEvenOdd fillShapeEvenOdd;
WgPipelineFillStroke fillStroke; WgPipelineFillStroke fillStroke;
// fill pipelines // fill pipelines
WgPipelineClipMask clipMask;
WgPipelineSolid solid[3]; WgPipelineSolid solid[3];
WgPipelineLinear linear[3]; WgPipelineLinear linear[3];
WgPipelineRadial radial[3]; WgPipelineRadial radial[3];
@ -189,7 +226,11 @@ struct WgPipelines
WgPipelineBlend computeBlendSolid; WgPipelineBlend computeBlendSolid;
WgPipelineBlend computeBlendGradient; WgPipelineBlend computeBlendGradient;
WgPipelineBlend computeBlendImage; WgPipelineBlend computeBlendImage;
WgPipelineComposeBlend computeComposeBlend; WgPipelineBlendMask computeBlendSolidMask;
WgPipelineBlendMask computeBlendGradientMask;
WgPipelineBlendMask computeBlendImageMask;
WgPipelineMaskCompose computeMaskCompose;
WgPipelineCompose computeCompose;
WgPipelineAntiAliasing computeAntiAliasing; WgPipelineAntiAliasing computeAntiAliasing;
void initialize(WgContext& context); void initialize(WgContext& context);

View file

@ -262,8 +262,17 @@ void WgRenderSettings::release(WgContext& context)
void WgRenderDataPaint::release(WgContext& context) void WgRenderDataPaint::release(WgContext& context)
{ {
bindGroupPaint.release(); bindGroupPaint.release();
clips.clear();
}; };
void WgRenderDataPaint::updateClips(tvg::Array<tvg::RenderData> &clips) {
this->clips.clear();
for (uint32_t i = 0; i < clips.count; i++)
if (clips[i])
this->clips.push((WgRenderDataPaint*)clips[i]);
}
//*********************************************************************** //***********************************************************************
// WgRenderDataShape // WgRenderDataShape
//*********************************************************************** //***********************************************************************
@ -371,6 +380,7 @@ void WgRenderDataShape::releaseMeshes(WgContext &context)
meshGroupShapes.release(context); meshGroupShapes.release(context);
pMin = {FLT_MAX, FLT_MAX}; pMin = {FLT_MAX, FLT_MAX};
pMax = {0.0f, 0.0f}; pMax = {0.0f, 0.0f};
clips.clear();
} }
@ -407,6 +417,7 @@ void WgRenderDataShapePool::free(WgContext& context, WgRenderDataShape* dataShap
dataShape->meshGroupShapesBBox.release(context); dataShape->meshGroupShapesBBox.release(context);
dataShape->meshGroupStrokes.release(context); dataShape->meshGroupStrokes.release(context);
dataShape->meshGroupStrokesBBox.release(context); dataShape->meshGroupStrokesBBox.release(context);
dataShape->clips.clear();
mPool.push(dataShape); mPool.push(dataShape);
} }

View file

@ -88,10 +88,13 @@ struct WgRenderDataPaint
WgBindGroupPaint bindGroupPaint{}; WgBindGroupPaint bindGroupPaint{};
RenderRegion viewport{}; RenderRegion viewport{};
float opacity{}; float opacity{};
Array<WgRenderDataPaint*> clips;
virtual ~WgRenderDataPaint() {}; virtual ~WgRenderDataPaint() {};
virtual void release(WgContext& context); virtual void release(WgContext& context);
virtual Type type() { return Type::Undefined; }; virtual Type type() { return Type::Undefined; };
void updateClips(tvg::Array<tvg::RenderData> &clips);
}; };
struct WgRenderDataShape: public WgRenderDataPaint struct WgRenderDataShape: public WgRenderDataPaint

View file

@ -85,7 +85,6 @@ void WgRenderStorage::renderShape(WgContext& context, WgRenderDataShape* renderD
{ {
assert(renderData); assert(renderData);
assert(mRenderPassEncoder); assert(mRenderPassEncoder);
// draw strokes
if (renderData->strokeFirst) { if (renderData->strokeFirst) {
drawStroke(context, renderData, blendType); drawStroke(context, renderData, blendType);
drawShape(context, renderData, blendType); drawShape(context, renderData, blendType);
@ -102,12 +101,21 @@ void WgRenderStorage::renderPicture(WgContext& context, WgRenderDataPicture* ren
assert(mRenderPassEncoder); assert(mRenderPassEncoder);
uint8_t blend = (uint8_t)blendType; uint8_t blend = (uint8_t)blendType;
auto& vp = renderData->viewport; auto& vp = renderData->viewport;
if ((vp.w <= 0) || (vp.h <= 0)) return;
wgpuRenderPassEncoderSetScissorRect(mRenderPassEncoder, vp.x * samples, vp.y * samples, vp.w * samples, vp.h * samples); wgpuRenderPassEncoderSetScissorRect(mRenderPassEncoder, vp.x * samples, vp.y * samples, vp.w * samples, vp.h * samples);
wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 0); wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 0);
mPipelines->image[blend].use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint, renderData->bindGroupPicture); mPipelines->image[blend].use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint, renderData->bindGroupPicture);
renderData->meshData.drawImage(context, mRenderPassEncoder); renderData->meshData.drawImage(context, mRenderPassEncoder);
} }
void WgRenderStorage::renderClipPath(WgContext& context, WgRenderDataPaint* renderData) {
assert(renderData);
assert(mRenderPassEncoder);
if (renderData->type() == Type::Shape)
drawShapeClipPath(context, (WgRenderDataShape*)renderData);
else if (renderData->type() == Type::Picture)
drawPictureClipPath(context, (WgRenderDataPicture*)renderData);
}
void WgRenderStorage::drawShape(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType) void WgRenderStorage::drawShape(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType)
{ {
@ -115,10 +123,12 @@ void WgRenderStorage::drawShape(WgContext& context, WgRenderDataShape* renderDat
assert(mRenderPassEncoder); assert(mRenderPassEncoder);
assert(renderData->meshGroupShapes.meshes.count == renderData->meshGroupShapesBBox.meshes.count); assert(renderData->meshGroupShapes.meshes.count == renderData->meshGroupShapesBBox.meshes.count);
if (renderData->renderSettingsShape.skip) return; if (renderData->renderSettingsShape.skip) return;
// apply viewport
auto& vp = renderData->viewport; auto& vp = renderData->viewport;
if ((vp.w <= 0) || (vp.h <= 0)) return;
wgpuRenderPassEncoderSetScissorRect(mRenderPassEncoder, vp.x * samples, vp.y * samples, vp.w * samples, vp.h * samples); wgpuRenderPassEncoderSetScissorRect(mRenderPassEncoder, vp.x * samples, vp.y * samples, vp.w * samples, vp.h * samples);
wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 0);
// setup fill rule // setup fill rule
wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 0);
if (renderData->fillRule == FillRule::Winding) if (renderData->fillRule == FillRule::Winding)
mPipelines->fillShapeWinding.use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint); mPipelines->fillShapeWinding.use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint);
else else
@ -145,7 +155,9 @@ void WgRenderStorage::drawStroke(WgContext& context, WgRenderDataShape* renderDa
assert(mRenderPassEncoder); assert(mRenderPassEncoder);
assert(renderData->meshGroupStrokes.meshes.count == renderData->meshGroupStrokesBBox.meshes.count); assert(renderData->meshGroupStrokes.meshes.count == renderData->meshGroupStrokesBBox.meshes.count);
if (renderData->renderSettingsStroke.skip) return; if (renderData->renderSettingsStroke.skip) return;
// apply viewport
auto& vp = renderData->viewport; auto& vp = renderData->viewport;
if ((vp.w <= 0) || (vp.h <= 0)) return;
wgpuRenderPassEncoderSetScissorRect(mRenderPassEncoder, vp.x * samples, vp.y * samples, vp.w * samples, vp.h * samples); wgpuRenderPassEncoderSetScissorRect(mRenderPassEncoder, vp.x * samples, vp.y * samples, vp.w * samples, vp.h * samples);
// draw stroke geometry // draw stroke geometry
uint8_t blend = (uint8_t)blendType; uint8_t blend = (uint8_t)blendType;
@ -169,6 +181,35 @@ void WgRenderStorage::drawStroke(WgContext& context, WgRenderDataShape* renderDa
} }
void WgRenderStorage::drawShapeClipPath(WgContext& context, WgRenderDataShape* renderData) {
assert(renderData);
assert(renderData->type() == Type::Shape);
assert(renderData->meshGroupShapes.meshes.count == renderData->meshGroupShapesBBox.meshes.count);
// draw shape geometry
wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 0);
// setup fill rule
if (renderData->fillRule == FillRule::Winding)
mPipelines->fillShapeWinding.use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint);
else
mPipelines->fillShapeEvenOdd.use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint);
// draw to stencil (first pass)
for (uint32_t i = 0; i < renderData->meshGroupShapes.meshes.count; i++)
renderData->meshGroupShapes.meshes[i]->drawFan(context, mRenderPassEncoder);
// fill shape geometry (second pass)
mPipelines->clipMask.use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint);
renderData->meshDataBBox.drawFan(context, mRenderPassEncoder);
}
void WgRenderStorage::drawPictureClipPath(WgContext& context, WgRenderDataPicture* renderData) {
assert(renderData);
assert(renderData->type() == Type::Picture);
assert(mRenderPassEncoder);
wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 0);
mPipelines->clipMask.use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint);
renderData->meshData.drawImage(context, mRenderPassEncoder);
}
void WgRenderStorage::clear(WGPUCommandEncoder commandEncoder) void WgRenderStorage::clear(WGPUCommandEncoder commandEncoder)
{ {
assert(commandEncoder); assert(commandEncoder);
@ -179,7 +220,14 @@ void WgRenderStorage::clear(WGPUCommandEncoder commandEncoder)
} }
void WgRenderStorage::blend(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc, WgPipelineBlend* pipeline, WgBindGroupBlendMethod* blendMethod, WgBindGroupOpacity* opacity) { void WgRenderStorage::blend(
WgContext& context,
WGPUCommandEncoder commandEncoder,
WgPipelineBlend* pipeline,
WgRenderStorage* targetSrc,
WgBindGroupBlendMethod* blendMethod,
WgBindGroupOpacity* opacity)
{
assert(commandEncoder); assert(commandEncoder);
assert(targetSrc); assert(targetSrc);
assert(pipeline); assert(pipeline);
@ -190,7 +238,46 @@ void WgRenderStorage::blend(WGPUCommandEncoder commandEncoder, WgRenderStorage*
} }
void WgRenderStorage::composeBlend( void WgRenderStorage::blendMask(
WgContext& context,
WGPUCommandEncoder commandEncoder,
WgPipelineBlendMask* pipeline,
WgRenderStorage* texMsk,
WgRenderStorage* texSrc,
WgBindGroupBlendMethod* blendMethod,
WgBindGroupOpacity* opacity)
{
assert(commandEncoder);
assert(texSrc);
assert(texMsk);
WgBindGroupTexComposeBlend composeBlend;
composeBlend.initialize(context.device, context.queue, texSrc->texViewColor, texMsk->texViewColor, texViewColor);
WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder);
pipeline->use(computePassEncoder, composeBlend, *blendMethod, *opacity);
dispatchWorkgroups(computePassEncoder);
endComputePass(computePassEncoder);
composeBlend.release();
};
void WgRenderStorage::maskCompose(
WgContext& context,
WGPUCommandEncoder commandEncoder,
WgRenderStorage* texMsk0)
{
assert(commandEncoder);
assert(texMsk0);
WgBindGroupTexMaskCompose maskCompose;
maskCompose.initialize(context.device, context.queue, texMsk0->texViewColor, texViewColor);
WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder);
mPipelines->computeMaskCompose.use(computePassEncoder, maskCompose);
dispatchWorkgroups(computePassEncoder);
endComputePass(computePassEncoder);
maskCompose.release();
}
void WgRenderStorage::compose(
WgContext& context, WgContext& context,
WGPUCommandEncoder commandEncoder, WGPUCommandEncoder commandEncoder,
WgRenderStorage* texSrc, WgRenderStorage* texSrc,
@ -205,7 +292,7 @@ void WgRenderStorage::composeBlend(
WgBindGroupTexComposeBlend composeBlend; WgBindGroupTexComposeBlend composeBlend;
composeBlend.initialize(context.device, context.queue, texSrc->texViewColor, texMsk->texViewColor, texViewColor); composeBlend.initialize(context.device, context.queue, texSrc->texViewColor, texMsk->texViewColor, texViewColor);
WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder); WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder);
mPipelines->computeComposeBlend.use(computePassEncoder, composeBlend, *composeMethod, *blendMethod, *opacity); mPipelines->computeCompose.use(computePassEncoder, composeBlend, *composeMethod, *blendMethod, *opacity);
dispatchWorkgroups(computePassEncoder); dispatchWorkgroups(computePassEncoder);
endComputePass(computePassEncoder); endComputePass(computePassEncoder);
composeBlend.release(); composeBlend.release();

View file

@ -52,10 +52,29 @@ public:
void renderShape(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType); void renderShape(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType);
void renderPicture(WgContext& context, WgRenderDataPicture* renderData, WgPipelineBlendType blendType); void renderPicture(WgContext& context, WgRenderDataPicture* renderData, WgPipelineBlendType blendType);
void renderClipPath(WgContext& context, WgRenderDataPaint* renderData);
void clear(WGPUCommandEncoder commandEncoder); void clear(WGPUCommandEncoder commandEncoder);
void blend(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc, WgPipelineBlend* pipeline, WgBindGroupBlendMethod* blendMethod, WgBindGroupOpacity* opacity); void blend(
void composeBlend( WgContext& context,
WGPUCommandEncoder commandEncoder,
WgPipelineBlend* pipeline,
WgRenderStorage* targetSrc,
WgBindGroupBlendMethod* blendMethod,
WgBindGroupOpacity* opacity);
void blendMask(
WgContext& context,
WGPUCommandEncoder commandEncoder,
WgPipelineBlendMask* pipeline,
WgRenderStorage* texMsk,
WgRenderStorage* texSrc,
WgBindGroupBlendMethod* blendMethod,
WgBindGroupOpacity* opacity);
void maskCompose(
WgContext& context,
WGPUCommandEncoder commandEncoder,
WgRenderStorage* texMsk0);
void compose(
WgContext& context, WgContext& context,
WGPUCommandEncoder commandEncoder, WGPUCommandEncoder commandEncoder,
WgRenderStorage* texMsk, WgRenderStorage* texMsk,
@ -68,6 +87,9 @@ private:
void drawShape(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType); void drawShape(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType);
void drawStroke(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType); void drawStroke(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType);
void drawShapeClipPath(WgContext& context, WgRenderDataShape* renderData);
void drawPictureClipPath(WgContext& context, WgRenderDataPicture* renderData);
void dispatchWorkgroups(WGPUComputePassEncoder computePassEncoder); void dispatchWorkgroups(WGPUComputePassEncoder computePassEncoder);
WGPUComputePassEncoder beginComputePass(WGPUCommandEncoder commandEncoder); WGPUComputePassEncoder beginComputePass(WGPUCommandEncoder commandEncoder);

View file

@ -67,6 +67,7 @@ void WgRenderer::release()
mBlendMethodPool.release(mContext); mBlendMethodPool.release(mContext);
mOpacityPool.release(mContext); mOpacityPool.release(mContext);
mRenderStorageRoot.release(mContext); mRenderStorageRoot.release(mContext);
mRenderStorageMask.release(mContext);
mRenderStorageScreen.release(mContext); mRenderStorageScreen.release(mContext);
mRenderStorageInterm.release(mContext); mRenderStorageInterm.release(mContext);
mPipelines.release(); mPipelines.release();
@ -119,6 +120,9 @@ RenderData WgRenderer::prepare(const RenderShape& rshape, RenderData data, const
if (rshape.stroke) if (rshape.stroke)
renderDataShape->renderSettingsStroke.update(mContext, rshape.stroke->fill, rshape.stroke->color, flags); renderDataShape->renderSettingsStroke.update(mContext, rshape.stroke->fill, rshape.stroke->color, flags);
// store clips data
renderDataShape->updateClips(clips);
return renderDataShape; return renderDataShape;
} }
@ -159,6 +163,9 @@ RenderData WgRenderer::prepare(Surface* surface, const RenderMesh* mesh, RenderD
renderDataPicture->imageData.textureView); renderDataPicture->imageData.textureView);
} }
// store clips data
renderDataPicture->updateClips(clips);
return renderDataPicture; return renderDataPicture;
} }
@ -175,6 +182,19 @@ bool WgRenderer::preRender()
} }
void WgRenderer::renderClipPath(Array<WgRenderDataPaint*>& clips)
{
for (uint32_t i = 0; i < clips.count; i++) {
renderClipPath(clips[i]->clips);
// render image to render target
mRenderStorageInterm.beginRenderPass(mCommandEncoder, true);
mRenderStorageInterm.renderClipPath(mContext, clips[i]);
mRenderStorageInterm.endRenderPass();
mRenderStorageMask.maskCompose(mContext, mCommandEncoder, &mRenderStorageInterm);
}
}
bool WgRenderer::renderShape(RenderData data) bool WgRenderer::renderShape(RenderData data)
{ {
// get current render storage // get current render storage
@ -183,11 +203,34 @@ bool WgRenderer::renderShape(RenderData data)
WgPipelineBlendType blendType = WgPipelines::blendMethodToBlendType(mBlendMethod); WgPipelineBlendType blendType = WgPipelines::blendMethodToBlendType(mBlendMethod);
WgRenderStorage* renderStorage = mRenderStorageStack.last(); WgRenderStorage* renderStorage = mRenderStorageStack.last();
assert(renderStorage); assert(renderStorage);
// use masked blend
if (dataShape->clips.count > 0) {
// terminate current render pass
renderStorage->endRenderPass();
// render clip path
mRenderStorageMask.beginRenderPass(mCommandEncoder, true);
mRenderStorageMask.renderClipPath(mContext, dataShape->clips[0]);
mRenderStorageMask.endRenderPass();
renderClipPath(dataShape->clips);
// render image to render target
mRenderStorageInterm.beginRenderPass(mCommandEncoder, true);
mRenderStorageInterm.renderShape(mContext, dataShape, blendType);
mRenderStorageInterm.endRenderPass();
// blend shape with current render storage
WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod);
WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, 255);
WgPipelineBlendMask* pipeline = &mContext.pipelines->computeBlendSolidMask;
if (dataShape->renderSettingsShape.fillType != WgRenderSettingsType::Solid)
pipeline = &mContext.pipelines->computeBlendGradientMask;
renderStorage->blendMask(mContext, mCommandEncoder,
pipeline, &mRenderStorageMask, &mRenderStorageInterm, blendMethod, opacity);
// restore current render pass
renderStorage->beginRenderPass(mCommandEncoder, false);
// use hardware blend // use hardware blend
if (WgPipelines::isBlendMethodSupportsHW(mBlendMethod)) } else if (WgPipelines::isBlendMethodSupportsHW(mBlendMethod)) {
renderStorage->renderShape(mContext, dataShape, blendType); renderStorage->renderShape(mContext, dataShape, blendType);
// use custom blend // use custom blend
else { } else {
// terminate current render pass // terminate current render pass
renderStorage->endRenderPass(); renderStorage->endRenderPass();
// render image to render target // render image to render target
@ -200,7 +243,8 @@ bool WgRenderer::renderShape(RenderData data)
WgPipelineBlend* pipeline = &mContext.pipelines->computeBlendSolid; WgPipelineBlend* pipeline = &mContext.pipelines->computeBlendSolid;
if (dataShape->renderSettingsShape.fillType != WgRenderSettingsType::Solid) if (dataShape->renderSettingsShape.fillType != WgRenderSettingsType::Solid)
pipeline = &mContext.pipelines->computeBlendGradient; pipeline = &mContext.pipelines->computeBlendGradient;
renderStorage->blend(mCommandEncoder, &mRenderStorageInterm, pipeline, blendMethod, opacity); renderStorage->blend(mContext, mCommandEncoder,
pipeline, &mRenderStorageInterm, blendMethod, opacity);
// restore current render pass // restore current render pass
renderStorage->beginRenderPass(mCommandEncoder, false); renderStorage->beginRenderPass(mCommandEncoder, false);
} }
@ -210,26 +254,50 @@ bool WgRenderer::renderShape(RenderData data)
bool WgRenderer::renderImage(RenderData data) bool WgRenderer::renderImage(RenderData data)
{ {
// get current render storage
WgRenderDataPicture *dataPicture = (WgRenderDataPicture*)data;
// get current render storage // get current render storage
WgPipelineBlendType blendType = WgPipelines::blendMethodToBlendType(mBlendMethod); WgPipelineBlendType blendType = WgPipelines::blendMethodToBlendType(mBlendMethod);
WgRenderStorage* renderStorage = mRenderStorageStack.last(); WgRenderStorage* renderStorage = mRenderStorageStack.last();
assert(renderStorage); assert(renderStorage);
// use masked blend
if (dataPicture->clips.count > 0) {
// terminate current render pass
renderStorage->endRenderPass();
// render clip path
mRenderStorageMask.beginRenderPass(mCommandEncoder, true);
mRenderStorageMask.renderClipPath(mContext, dataPicture->clips[0]);
mRenderStorageMask.endRenderPass();
renderClipPath(dataPicture->clips);
// render image to render target
mRenderStorageInterm.beginRenderPass(mCommandEncoder, true);
mRenderStorageInterm.renderPicture(mContext, dataPicture, blendType);
mRenderStorageInterm.endRenderPass();
// blend shape with current render storage
WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod);
WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, 255);
WgPipelineBlendMask* pipeline = &mContext.pipelines->computeBlendImageMask;
renderStorage->blendMask(mContext, mCommandEncoder,
pipeline, &mRenderStorageMask, &mRenderStorageInterm, blendMethod, opacity);
// restore current render pass
renderStorage->beginRenderPass(mCommandEncoder, false);
// use hardware blend // use hardware blend
if (WgPipelines::isBlendMethodSupportsHW(mBlendMethod)) } else if (WgPipelines::isBlendMethodSupportsHW(mBlendMethod))
renderStorage->renderPicture(mContext, (WgRenderDataPicture *)data, blendType); renderStorage->renderPicture(mContext, dataPicture, blendType);
// use custom blend // use custom blend
else { else {
// terminate current render pass // terminate current render pass
renderStorage->endRenderPass(); renderStorage->endRenderPass();
// render image to render target // render image to render target
mRenderStorageInterm.beginRenderPass(mCommandEncoder, true); mRenderStorageInterm.beginRenderPass(mCommandEncoder, true);
mRenderStorageInterm.renderPicture(mContext, (WgRenderDataPicture *)data, blendType); mRenderStorageInterm.renderPicture(mContext, dataPicture, blendType);
mRenderStorageInterm.endRenderPass(); mRenderStorageInterm.endRenderPass();
// blend shape with current render storage // blend shape with current render storage
WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod); WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod);
WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, 255); WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, 255);
WgPipelineBlend* pipeline = &mContext.pipelines->computeBlendImage; WgPipelineBlend* pipeline = &mContext.pipelines->computeBlendImage;
renderStorage->blend(mCommandEncoder, &mRenderStorageInterm, pipeline, blendMethod, opacity); renderStorage->blend(mContext, mCommandEncoder,
pipeline, &mRenderStorageInterm, blendMethod, opacity);
// restore current render pass // restore current render pass
renderStorage->beginRenderPass(mCommandEncoder, false); renderStorage->beginRenderPass(mCommandEncoder, false);
} }
@ -258,7 +326,7 @@ void WgRenderer::dispose(RenderData data) {
RenderRegion WgRenderer::region(TVG_UNUSED RenderData data) RenderRegion WgRenderer::region(TVG_UNUSED RenderData data)
{ {
return { 0, 0, INT32_MAX, INT32_MAX }; return { 0, 0, (int32_t)mTargetSurface.w, (int32_t)mTargetSurface.h };
} }
@ -355,6 +423,7 @@ bool WgRenderer::target(WGPUInstance instance, WGPUSurface surface, uint32_t w,
initialize(); initialize();
mRenderStorageInterm.initialize(mContext, w, h, WG_SSAA_SAMPLES); mRenderStorageInterm.initialize(mContext, w, h, WG_SSAA_SAMPLES);
mRenderStorageMask.initialize(mContext, w, h, WG_SSAA_SAMPLES);
mRenderStorageRoot.initialize(mContext, w, h, WG_SSAA_SAMPLES); mRenderStorageRoot.initialize(mContext, w, h, WG_SSAA_SAMPLES);
mRenderStorageScreen.initialize(mContext, w, h, 1, WGPUTextureFormat_BGRA8Unorm); mRenderStorageScreen.initialize(mContext, w, h, 1, WGPUTextureFormat_BGRA8Unorm);
return true; return true;
@ -400,7 +469,9 @@ bool WgRenderer::endComposite(TVG_UNUSED Compositor* cmp)
// blent scene to current render storage // blent scene to current render storage
WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, comp->blendMethod); WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, comp->blendMethod);
WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, comp->opacity); WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, comp->opacity);
mRenderStorageStack.last()->blend(mCommandEncoder, renderStorageSrc, &mContext.pipelines->computeBlendImage, blendMethod, opacity); mRenderStorageStack.last()->blend(mContext, mCommandEncoder,
&mContext.pipelines->computeBlendImage,
renderStorageSrc, blendMethod, opacity);
// back render targets to the pool // back render targets to the pool
mRenderStoragePool.free(mContext, renderStorageSrc); mRenderStoragePool.free(mContext, renderStorageSrc);
@ -424,7 +495,7 @@ bool WgRenderer::endComposite(TVG_UNUSED Compositor* cmp)
// compose and blend // compose and blend
// dest = blend(dest, compose(src, msk, composeMethod), blendMethod, opacity) // dest = blend(dest, compose(src, msk, composeMethod), blendMethod, opacity)
renderStorageDst->composeBlend(mContext, mCommandEncoder, renderStorageDst->compose(mContext, mCommandEncoder,
renderStorageSrc, renderStorageMsk, renderStorageSrc, renderStorageMsk,
composeMethod, blendMethod, opacity); composeMethod, blendMethod, opacity);

View file

@ -62,9 +62,11 @@ private:
void initialize(); void initialize();
void release(); void release();
void clearDisposes(); void clearDisposes();
void renderClipPath(Array<WgRenderDataPaint*>& clips);
WGPUCommandEncoder mCommandEncoder{}; // render handles WGPUCommandEncoder mCommandEncoder{};
WgRenderStorage mRenderStorageInterm; // intermidiate buffer to render WgRenderStorage mRenderStorageInterm; // intermidiate buffer to render
WgRenderStorage mRenderStorageMask; // buffer to render mask
WgRenderStorage mRenderStorageRoot; // root render storage WgRenderStorage mRenderStorageRoot; // root render storage
WgRenderStorage mRenderStorageScreen; // storage with data after antializing WgRenderStorage mRenderStorageScreen; // storage with data after antializing
WgRenderStoragePool mRenderStoragePool; // pool to hold render tree storages WgRenderStoragePool mRenderStoragePool; // pool to hold render tree storages

View file

@ -53,8 +53,8 @@ fn vs_main(in: VertexInput) -> VertexOutput {
} }
@fragment @fragment
fn fs_main(in: VertexOutput) -> void { fn fs_main(in: VertexOutput) -> @location(0) vec4f {
// nothing to draw, just stencil value return vec4f(0.0, 0.0, 0.0, 1.0);
} }
); );
@ -393,6 +393,34 @@ fn cs_main( @builtin(global_invocation_id) id: vec3u) {
var Ra: f32 = 1.0; var Ra: f32 = 1.0;
); );
const std::string strBlendMaskShaderHeader = WG_SHADER_SOURCE(
@group(0) @binding(0) var imageSrc : texture_storage_2d<rgba8unorm, read>;
@group(0) @binding(1) var imageMsk : texture_storage_2d<rgba8unorm, read>;
@group(0) @binding(2) var imageDst : texture_storage_2d<rgba8unorm, read_write>;
@group(1) @binding(0) var<uniform> blendMethod : u32;
@group(2) @binding(0) var<uniform> opacity : f32;
@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; };
let colorMsk = textureLoad(imageMsk, id.xy);
if (colorMsk.a == 0.0) { return; }
let colorSrc = textureLoad(imageSrc, id.xy);
if (colorSrc.a == 0.0) { return; }
let colorDst = textureLoad(imageDst, id.xy);
var One: vec3f = vec3(1.0);
var So: f32 = opacity;
var Sc: vec3f = colorSrc.rgb;
var Sa: f32 = colorSrc.a;
var Dc: vec3f = colorDst.rgb;
var Da: f32 = colorDst.a;
var Rc: vec3f = colorDst.rgb;
var Ra: f32 = 1.0;
);
const std::string strBlendShaderPreConditionsGradient = WG_SHADER_SOURCE( const std::string strBlendShaderPreConditionsGradient = WG_SHADER_SOURCE(
Sc = Sc + Dc.rgb * (1.0 - Sa); Sc = Sc + Dc.rgb * (1.0 - Sa);
Sa = Sa + Da * (1.0 - Sa); Sa = Sa + Da * (1.0 - Sa);
@ -484,6 +512,31 @@ const std::string strComputeBlendImage =
strBlendShaderFooter; strBlendShaderFooter;
const char* cShaderSource_PipelineComputeBlendImage = strComputeBlendImage.c_str(); const char* cShaderSource_PipelineComputeBlendImage = strComputeBlendImage.c_str();
// pipeline shader modules blend solid mask
const std::string strComputeBlendSolidMask =
strBlendMaskShaderHeader +
strBlendShaderBlendMethod +
strBlendShaderFooter;
const char* cShaderSource_PipelineComputeBlendSolidMask = strComputeBlendSolidMask.c_str();
// pipeline shader modules blend gradient mask
const std::string strComputeBlendGradientMask =
strBlendMaskShaderHeader +
strBlendShaderPreConditionsGradient +
strBlendShaderBlendMethod +
strBlendShaderPostConditionsGradient +
strBlendShaderFooter;
const char* cShaderSource_PipelineComputeBlendGradientMask = strComputeBlendGradientMask.c_str();
// pipeline shader modules blend image mask
const std::string strComputeBlendImageMask =
strBlendMaskShaderHeader +
strBlendShaderPreConditionsImage +
strBlendShaderBlendMethod +
strBlendShaderPostConditionsImage +
strBlendShaderFooter;
const char* cShaderSource_PipelineComputeBlendImageMask = strComputeBlendImageMask.c_str();
// pipeline shader modules clear // pipeline shader modules clear
const char* cShaderSource_PipelineComputeClear = WG_SHADER_SOURCE( const char* cShaderSource_PipelineComputeClear = WG_SHADER_SOURCE(
@group(0) @binding(0) var imageDst : texture_storage_2d<rgba8unorm, read_write>; @group(0) @binding(0) var imageDst : texture_storage_2d<rgba8unorm, read_write>;
@ -494,8 +547,21 @@ fn cs_main( @builtin(global_invocation_id) id: vec3u) {
} }
); );
// pipeline shader modules compose blend // pipeline shader modules compose
const char* cShaderSource_PipelineComputeComposeBlend = WG_SHADER_SOURCE( const char* cShaderSource_PipelineComputeMaskCompose = WG_SHADER_SOURCE(
@group(0) @binding(0) var imageMsk0 : texture_storage_2d<rgba8unorm, read>;
@group(0) @binding(1) var imageMsk1 : texture_storage_2d<rgba8unorm, read_write>;
@compute @workgroup_size(8, 8)
fn cs_main( @builtin(global_invocation_id) id: vec3u) {
let colorMsk0 = textureLoad(imageMsk0, id.xy);
let colorMsk1 = textureLoad(imageMsk1, id.xy);
textureStore(imageMsk1, id.xy, colorMsk0 * colorMsk1);
}
);
// pipeline shader modules compose
const char* cShaderSource_PipelineComputeCompose = WG_SHADER_SOURCE(
@group(0) @binding(0) var imageSrc : texture_storage_2d<rgba8unorm, read>; @group(0) @binding(0) var imageSrc : texture_storage_2d<rgba8unorm, read>;
@group(0) @binding(1) var imageMsk : texture_storage_2d<rgba8unorm, read>; @group(0) @binding(1) var imageMsk : texture_storage_2d<rgba8unorm, read>;
@group(0) @binding(2) var imageDst : texture_storage_2d<rgba8unorm, read_write>; @group(0) @binding(2) var imageDst : texture_storage_2d<rgba8unorm, read_write>;

View file

@ -45,7 +45,11 @@ extern const char* cShaderSource_PipelineComputeClear;
extern const char* cShaderSource_PipelineComputeBlendSolid; extern const char* cShaderSource_PipelineComputeBlendSolid;
extern const char* cShaderSource_PipelineComputeBlendGradient; extern const char* cShaderSource_PipelineComputeBlendGradient;
extern const char* cShaderSource_PipelineComputeBlendImage; extern const char* cShaderSource_PipelineComputeBlendImage;
extern const char* cShaderSource_PipelineComputeComposeBlend; extern const char* cShaderSource_PipelineComputeBlendSolidMask;
extern const char* cShaderSource_PipelineComputeBlendGradientMask;
extern const char* cShaderSource_PipelineComputeBlendImageMask;
extern const char* cShaderSource_PipelineComputeMaskCompose;
extern const char* cShaderSource_PipelineComputeCompose;
extern const char* cShaderSource_PipelineComputeAntiAlias; extern const char* cShaderSource_PipelineComputeAntiAlias;
#endif // _TVG_WG_SHADER_SRC_H_ #endif // _TVG_WG_SHADER_SRC_H_