mirror of
https://github.com/thorvg/thorvg.git
synced 2025-06-14 12:04:29 +00:00
wg_engine: fix blend methods support
Full review of blending support. Support Solid color, Gradient fill and Image blending workflows See Blending, SceneBlending, Opacity examples
This commit is contained in:
parent
35d168db9e
commit
e27ce23d4b
10 changed files with 210 additions and 144 deletions
|
@ -596,36 +596,21 @@ void WgRenderPipeline::set(WGPURenderPassEncoder renderPassEncoder)
|
||||||
WGPUBlendState WgRenderPipeline::makeBlendState(WgPipelineBlendType blendType)
|
WGPUBlendState WgRenderPipeline::makeBlendState(WgPipelineBlendType blendType)
|
||||||
{
|
{
|
||||||
WGPUBlendState blendState{};
|
WGPUBlendState blendState{};
|
||||||
// src
|
// srcover
|
||||||
if (blendType == WgPipelineBlendType::Src) {
|
if (blendType == WgPipelineBlendType::SrcOver) {
|
||||||
blendState.color.operation = WGPUBlendOperation_Add;
|
blendState.color.operation = WGPUBlendOperation_Add;
|
||||||
blendState.color.srcFactor = WGPUBlendFactor_One;
|
blendState.color.srcFactor = WGPUBlendFactor_One;
|
||||||
blendState.color.dstFactor = WGPUBlendFactor_Zero;
|
blendState.color.dstFactor = WGPUBlendFactor_Zero;
|
||||||
} else // normal
|
// normal
|
||||||
if (blendType == WgPipelineBlendType::Normal) {
|
} else if (blendType == WgPipelineBlendType::Normal) {
|
||||||
blendState.color.operation = WGPUBlendOperation_Add;
|
blendState.color.operation = WGPUBlendOperation_Add;
|
||||||
blendState.color.srcFactor = WGPUBlendFactor_One;
|
blendState.color.srcFactor = WGPUBlendFactor_One;
|
||||||
blendState.color.dstFactor = WGPUBlendFactor_OneMinusSrcAlpha;
|
blendState.color.dstFactor = WGPUBlendFactor_OneMinusSrcAlpha;
|
||||||
} else // add
|
// custom
|
||||||
if (blendType == WgPipelineBlendType::Add) {
|
} else if (blendType == WgPipelineBlendType::Custom) {
|
||||||
blendState.color.operation = WGPUBlendOperation_Add;
|
blendState.color.operation = WGPUBlendOperation_Add;
|
||||||
blendState.color.srcFactor = WGPUBlendFactor_One;
|
blendState.color.srcFactor = WGPUBlendFactor_One;
|
||||||
blendState.color.dstFactor = WGPUBlendFactor_One;
|
|
||||||
} else // mult
|
|
||||||
if (blendType == WgPipelineBlendType::Mult) {
|
|
||||||
blendState.color.operation = WGPUBlendOperation_Add;
|
|
||||||
blendState.color.srcFactor = WGPUBlendFactor_Dst;
|
|
||||||
blendState.color.dstFactor = WGPUBlendFactor_Zero;
|
blendState.color.dstFactor = WGPUBlendFactor_Zero;
|
||||||
} else // min
|
|
||||||
if (blendType == WgPipelineBlendType::Min) {
|
|
||||||
blendState.color.operation = WGPUBlendOperation_Min;
|
|
||||||
blendState.color.srcFactor = WGPUBlendFactor_One;
|
|
||||||
blendState.color.dstFactor = WGPUBlendFactor_One;
|
|
||||||
} else // max
|
|
||||||
if (blendType == WgPipelineBlendType::Max) {
|
|
||||||
blendState.color.operation = WGPUBlendOperation_Max;
|
|
||||||
blendState.color.srcFactor = WGPUBlendFactor_One;
|
|
||||||
blendState.color.dstFactor = WGPUBlendFactor_One;
|
|
||||||
}
|
}
|
||||||
blendState.alpha = blendState.color;
|
blendState.alpha = blendState.color;
|
||||||
return blendState;
|
return blendState;
|
||||||
|
|
|
@ -36,12 +36,9 @@ struct WgCompositor: public Compositor {
|
||||||
};
|
};
|
||||||
|
|
||||||
enum class WgPipelineBlendType {
|
enum class WgPipelineBlendType {
|
||||||
Src = 0, // S
|
SrcOver = 0,
|
||||||
Normal, // (Sa * S) + (255 - Sa) * D
|
Normal,
|
||||||
Add, // (S + D)
|
Custom
|
||||||
Mult, // (S * D)
|
|
||||||
Min, // min(S, D)
|
|
||||||
Max // max(S, D)
|
|
||||||
};
|
};
|
||||||
|
|
||||||
struct WgPipelines;
|
struct WgPipelines;
|
||||||
|
|
|
@ -55,7 +55,7 @@ void WgPipelineFillShapeWinding::initialize(WGPUDevice device)
|
||||||
auto pipelineLabel = "The render pipeline fill shape winding";
|
auto pipelineLabel = "The render pipeline fill shape winding";
|
||||||
|
|
||||||
// allocate all pipeline handles
|
// allocate all pipeline handles
|
||||||
allocate(device, WgPipelineBlendType::Add, WGPUColorWriteMask_None,
|
allocate(device, WgPipelineBlendType::SrcOver, WGPUColorWriteMask_None,
|
||||||
vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts),
|
vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts),
|
||||||
bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts),
|
bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts),
|
||||||
stencilFunctionFront, stencilOperationFront, stencilFunctionBack, stencilOperationBack,
|
stencilFunctionFront, stencilOperationFront, stencilFunctionBack, stencilOperationBack,
|
||||||
|
@ -89,7 +89,7 @@ void WgPipelineFillShapeEvenOdd::initialize(WGPUDevice device)
|
||||||
auto pipelineLabel = "The render pipeline fill shape Even Odd";
|
auto pipelineLabel = "The render pipeline fill shape Even Odd";
|
||||||
|
|
||||||
// allocate all pipeline handles
|
// allocate all pipeline handles
|
||||||
allocate(device, WgPipelineBlendType::Add, WGPUColorWriteMask_None,
|
allocate(device, WgPipelineBlendType::SrcOver, WGPUColorWriteMask_None,
|
||||||
vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts),
|
vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts),
|
||||||
bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts),
|
bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts),
|
||||||
stencilFunctionFront, stencilOperationFront, stencilFunctionBack, stencilOperationBack,
|
stencilFunctionFront, stencilOperationFront, stencilFunctionBack, stencilOperationBack,
|
||||||
|
@ -121,7 +121,7 @@ void WgPipelineFillStroke::initialize(WGPUDevice device)
|
||||||
auto pipelineLabel = "The render pipeline fill stroke";
|
auto pipelineLabel = "The render pipeline fill stroke";
|
||||||
|
|
||||||
// allocate all pipeline handles
|
// allocate all pipeline handles
|
||||||
allocate(device, WgPipelineBlendType::Add, WGPUColorWriteMask_None,
|
allocate(device, WgPipelineBlendType::SrcOver, WGPUColorWriteMask_None,
|
||||||
vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts),
|
vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts),
|
||||||
bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts),
|
bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts),
|
||||||
stencilFunction, stencilOperation, stencilFunction, stencilOperation,
|
stencilFunction, stencilOperation, stencilFunction, stencilOperation,
|
||||||
|
@ -285,7 +285,7 @@ void WgPipelineClear::initialize(WGPUDevice device)
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
void WgPipelineBlend::initialize(WGPUDevice device)
|
void WgPipelineBlend::initialize(WGPUDevice device, const char *shaderSource)
|
||||||
{
|
{
|
||||||
// bind groups and layouts
|
// bind groups and layouts
|
||||||
WGPUBindGroupLayout bindGroupLayouts[] = {
|
WGPUBindGroupLayout bindGroupLayouts[] = {
|
||||||
|
@ -296,7 +296,6 @@ void WgPipelineBlend::initialize(WGPUDevice device)
|
||||||
};
|
};
|
||||||
|
|
||||||
// sheder source and labels
|
// sheder source and labels
|
||||||
auto shaderSource = cShaderSource_PipelineComputeBlend;
|
|
||||||
auto shaderLabel = "The compute shader blend";
|
auto shaderLabel = "The compute shader blend";
|
||||||
auto pipelineLabel = "The compute pipeline blend";
|
auto pipelineLabel = "The compute pipeline blend";
|
||||||
|
|
||||||
|
@ -380,7 +379,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);
|
||||||
for (uint8_t type = (uint8_t)WgPipelineBlendType::Src; type <= (uint8_t)WgPipelineBlendType::Max; 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);
|
||||||
radial[type].initialize(context.device, (WgPipelineBlendType)type);
|
radial[type].initialize(context.device, (WgPipelineBlendType)type);
|
||||||
|
@ -388,7 +387,9 @@ void WgPipelines::initialize(WgContext& context)
|
||||||
}
|
}
|
||||||
// compute pipelines
|
// compute pipelines
|
||||||
computeClear.initialize(context.device);
|
computeClear.initialize(context.device);
|
||||||
computeBlend.initialize(context.device);
|
computeBlendSolid.initialize(context.device, cShaderSource_PipelineComputeBlendSolid);
|
||||||
|
computeBlendGradient.initialize(context.device, cShaderSource_PipelineComputeBlendGradient);
|
||||||
|
computeBlendImage.initialize(context.device, cShaderSource_PipelineComputeBlendImage);
|
||||||
computeCompose.initialize(context.device);
|
computeCompose.initialize(context.device);
|
||||||
computeComposeBlend.initialize(context.device);
|
computeComposeBlend.initialize(context.device);
|
||||||
computeAntiAliasing.initialize(context.device);
|
computeAntiAliasing.initialize(context.device);
|
||||||
|
@ -415,10 +416,12 @@ void WgPipelines::release()
|
||||||
computeAntiAliasing.release();
|
computeAntiAliasing.release();
|
||||||
computeComposeBlend.release();
|
computeComposeBlend.release();
|
||||||
computeCompose.release();
|
computeCompose.release();
|
||||||
computeBlend.release();
|
computeBlendImage.release();
|
||||||
|
computeBlendGradient.release();
|
||||||
|
computeBlendSolid.release();
|
||||||
computeClear.release();
|
computeClear.release();
|
||||||
// fill pipelines
|
// fill pipelines
|
||||||
for (uint8_t type = (uint8_t)WgPipelineBlendType::Src; type <= (uint8_t)WgPipelineBlendType::Max; type++) {
|
for (uint8_t type = (uint8_t)WgPipelineBlendType::SrcOver; type <= (uint8_t)WgPipelineBlendType::Custom; type++) {
|
||||||
image[type].release();
|
image[type].release();
|
||||||
radial[type].release();
|
radial[type].release();
|
||||||
linear[type].release();
|
linear[type].release();
|
||||||
|
@ -435,10 +438,6 @@ bool WgPipelines::isBlendMethodSupportsHW(BlendMethod blendMethod)
|
||||||
switch (blendMethod) {
|
switch (blendMethod) {
|
||||||
case BlendMethod::SrcOver:
|
case BlendMethod::SrcOver:
|
||||||
case BlendMethod::Normal:
|
case BlendMethod::Normal:
|
||||||
case BlendMethod::Add:
|
|
||||||
case BlendMethod::Multiply:
|
|
||||||
case BlendMethod::Darken:
|
|
||||||
case BlendMethod::Lighten:
|
|
||||||
return true;
|
return true;
|
||||||
default: return false;
|
default: return false;
|
||||||
};
|
};
|
||||||
|
@ -448,12 +447,8 @@ bool WgPipelines::isBlendMethodSupportsHW(BlendMethod blendMethod)
|
||||||
WgPipelineBlendType WgPipelines::blendMethodToBlendType(BlendMethod blendMethod)
|
WgPipelineBlendType WgPipelines::blendMethodToBlendType(BlendMethod blendMethod)
|
||||||
{
|
{
|
||||||
switch (blendMethod) {
|
switch (blendMethod) {
|
||||||
case BlendMethod::SrcOver: return WgPipelineBlendType::Src;
|
case BlendMethod::SrcOver: return WgPipelineBlendType::SrcOver;
|
||||||
case BlendMethod::Normal: return WgPipelineBlendType::Normal;
|
case BlendMethod::Normal: return WgPipelineBlendType::Normal;
|
||||||
case BlendMethod::Add: return WgPipelineBlendType::Add;
|
default: return WgPipelineBlendType::Custom;
|
||||||
case BlendMethod::Multiply: return WgPipelineBlendType::Mult;
|
|
||||||
case BlendMethod::Darken: return WgPipelineBlendType::Min;
|
|
||||||
case BlendMethod::Lighten: return WgPipelineBlendType::Max;
|
|
||||||
default: return WgPipelineBlendType::Src;
|
|
||||||
};
|
};
|
||||||
}
|
}
|
||||||
|
|
|
@ -103,7 +103,7 @@ struct WgPipelineRadial: public WgRenderPipeline
|
||||||
|
|
||||||
struct WgPipelineImage: public WgRenderPipeline
|
struct WgPipelineImage: public WgRenderPipeline
|
||||||
{
|
{
|
||||||
void initialize(WGPUDevice device) override {}
|
void initialize(WGPUDevice device) override { assert(false); };
|
||||||
void initialize(WGPUDevice device, WgPipelineBlendType blendType);
|
void initialize(WGPUDevice device, WgPipelineBlendType blendType);
|
||||||
void use(WGPURenderPassEncoder encoder, WgBindGroupCanvas& groupCanvas, WgBindGroupPaint& groupPaint, WgBindGroupPicture& groupPicture)
|
void use(WGPURenderPassEncoder encoder, WgBindGroupCanvas& groupCanvas, WgBindGroupPaint& groupPaint, WgBindGroupPicture& groupPicture)
|
||||||
{
|
{
|
||||||
|
@ -131,7 +131,8 @@ struct WgPipelineClear: public WgComputePipeline
|
||||||
|
|
||||||
struct WgPipelineBlend: public WgComputePipeline
|
struct WgPipelineBlend: public WgComputePipeline
|
||||||
{
|
{
|
||||||
void initialize(WGPUDevice device) override;
|
void initialize(WGPUDevice device) override { assert(false); };
|
||||||
|
void initialize(WGPUDevice device, const char *shaderSource);
|
||||||
void use(WGPUComputePassEncoder encoder, WgBindGroupTextureStorageRgba& groupTexSrc, WgBindGroupTextureStorageRgba& groupTexDst, WgBindGroupBlendMethod& blendMethod, WgBindGroupOpacity& groupOpacity)
|
void use(WGPUComputePassEncoder encoder, WgBindGroupTextureStorageRgba& groupTexSrc, WgBindGroupTextureStorageRgba& groupTexDst, WgBindGroupBlendMethod& blendMethod, WgBindGroupOpacity& groupOpacity)
|
||||||
{
|
{
|
||||||
set(encoder);
|
set(encoder);
|
||||||
|
@ -193,13 +194,15 @@ struct WgPipelines
|
||||||
WgPipelineFillShapeEvenOdd fillShapeEvenOdd;
|
WgPipelineFillShapeEvenOdd fillShapeEvenOdd;
|
||||||
WgPipelineFillStroke fillStroke;
|
WgPipelineFillStroke fillStroke;
|
||||||
// fill pipelines
|
// fill pipelines
|
||||||
WgPipelineSolid solid[6];
|
WgPipelineSolid solid[3];
|
||||||
WgPipelineLinear linear[6];
|
WgPipelineLinear linear[3];
|
||||||
WgPipelineRadial radial[6];
|
WgPipelineRadial radial[3];
|
||||||
WgPipelineImage image[6];
|
WgPipelineImage image[3];
|
||||||
// compute pipelines
|
// compute pipelines
|
||||||
WgPipelineClear computeClear;
|
WgPipelineClear computeClear;
|
||||||
WgPipelineBlend computeBlend;
|
WgPipelineBlend computeBlendSolid;
|
||||||
|
WgPipelineBlend computeBlendGradient;
|
||||||
|
WgPipelineBlend computeBlendImage;
|
||||||
WgPipelineCompose computeCompose;
|
WgPipelineCompose computeCompose;
|
||||||
WgPipelineComposeBlend computeComposeBlend;
|
WgPipelineComposeBlend computeComposeBlend;
|
||||||
WgPipelineAntiAliasing computeAntiAliasing;
|
WgPipelineAntiAliasing computeAntiAliasing;
|
||||||
|
|
|
@ -173,15 +173,15 @@ void WgRenderStorage::clear(WGPUCommandEncoder commandEncoder)
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
void WgRenderStorage::blend(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc, WgBindGroupBlendMethod* blendMethod, WgBindGroupOpacity* opacity)
|
void WgRenderStorage::blend(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc, WgPipelineBlend* pipeline, WgBindGroupBlendMethod* blendMethod, WgBindGroupOpacity* opacity) {
|
||||||
{
|
|
||||||
assert(commandEncoder);
|
assert(commandEncoder);
|
||||||
assert(targetSrc);
|
assert(targetSrc);
|
||||||
|
assert(pipeline);
|
||||||
WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder);
|
WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder);
|
||||||
mPipelines->computeBlend.use(computePassEncoder, targetSrc->bindGroupTexStorageRgba, bindGroupTexStorageRgba, *blendMethod, *opacity);
|
pipeline->use(computePassEncoder, targetSrc->bindGroupTexStorageRgba, bindGroupTexStorageRgba, *blendMethod, *opacity);
|
||||||
dispatchWorkgroups(computePassEncoder);
|
dispatchWorkgroups(computePassEncoder);
|
||||||
endComputePass(computePassEncoder);
|
endComputePass(computePassEncoder);
|
||||||
};
|
}
|
||||||
|
|
||||||
|
|
||||||
void WgRenderStorage::compose(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetMsk, WgBindGroupCompositeMethod* composeMethod, WgBindGroupOpacity* opacity)
|
void WgRenderStorage::compose(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetMsk, WgBindGroupCompositeMethod* composeMethod, WgBindGroupOpacity* opacity)
|
||||||
|
|
|
@ -53,7 +53,7 @@ public:
|
||||||
void renderPicture(WgContext& context, WgRenderDataPicture* renderData, WgPipelineBlendType blendType);
|
void renderPicture(WgContext& context, WgRenderDataPicture* renderData, WgPipelineBlendType blendType);
|
||||||
|
|
||||||
void clear(WGPUCommandEncoder commandEncoder);
|
void clear(WGPUCommandEncoder commandEncoder);
|
||||||
void blend(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc, WgBindGroupBlendMethod* blendMethod, WgBindGroupOpacity* opacity);
|
void blend(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc, WgPipelineBlend* pipeline, WgBindGroupBlendMethod* blendMethod, WgBindGroupOpacity* opacity);
|
||||||
void compose(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetMsk, WgBindGroupCompositeMethod* composeMethod, WgBindGroupOpacity* opacity);
|
void compose(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetMsk, WgBindGroupCompositeMethod* composeMethod, WgBindGroupOpacity* opacity);
|
||||||
void composeBlend(
|
void composeBlend(
|
||||||
WgContext& context,
|
WgContext& context,
|
||||||
|
|
|
@ -63,7 +63,7 @@ void WgRenderer::release()
|
||||||
mOpacityPool.release(mContext);
|
mOpacityPool.release(mContext);
|
||||||
mRenderStorageRoot.release(mContext);
|
mRenderStorageRoot.release(mContext);
|
||||||
mRenderStorageScreen.release(mContext);
|
mRenderStorageScreen.release(mContext);
|
||||||
mRenderTarget.release(mContext);
|
mRenderStorageInterm.release(mContext);
|
||||||
mPipelines.release();
|
mPipelines.release();
|
||||||
}
|
}
|
||||||
|
|
||||||
|
@ -152,24 +152,29 @@ bool WgRenderer::preRender()
|
||||||
bool WgRenderer::renderShape(RenderData data)
|
bool WgRenderer::renderShape(RenderData data)
|
||||||
{
|
{
|
||||||
// get current render storage
|
// get current render storage
|
||||||
if (((WgRenderDataShape*)data)->opacity == 0) return 0;
|
WgRenderDataShape *dataShape = (WgRenderDataShape*)data;
|
||||||
|
if (dataShape->opacity == 0) return 0;
|
||||||
WgPipelineBlendType blendType = WgPipelines::blendMethodToBlendType(mBlendMethod);
|
WgPipelineBlendType blendType = WgPipelines::blendMethodToBlendType(mBlendMethod);
|
||||||
WgRenderStorage* renderStorage = mRenderStorageStack.last();
|
WgRenderStorage* renderStorage = mRenderStorageStack.last();
|
||||||
assert(renderStorage);
|
assert(renderStorage);
|
||||||
// use hardware blend
|
// use hardware blend
|
||||||
if (WgPipelines::isBlendMethodSupportsHW(mBlendMethod))
|
if (WgPipelines::isBlendMethodSupportsHW(mBlendMethod))
|
||||||
renderStorage->renderShape(mContext, (WgRenderDataShape *)data, blendType);
|
renderStorage->renderShape(mContext, dataShape, blendType);
|
||||||
else { // use custom blend
|
// use custom blend
|
||||||
|
else {
|
||||||
// terminate current render pass
|
// terminate current render pass
|
||||||
renderStorage->endRenderPass();
|
renderStorage->endRenderPass();
|
||||||
// render image to render target
|
// render image to render target
|
||||||
mRenderTarget.beginRenderPass(mCommandEncoder, true);
|
mRenderStorageInterm.beginRenderPass(mCommandEncoder, true);
|
||||||
mRenderTarget.renderShape(mContext, (WgRenderDataShape *)data, blendType);
|
mRenderStorageInterm.renderShape(mContext, dataShape, blendType);
|
||||||
mRenderTarget.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);
|
||||||
renderStorage->blend(mCommandEncoder, &mRenderTarget, blendMethod, opacity);
|
WgPipelineBlend* pipeline = &mContext.pipelines->computeBlendSolid;
|
||||||
|
if (dataShape->renderSettingsShape.fillType != WgRenderSettingsType::Solid)
|
||||||
|
pipeline = &mContext.pipelines->computeBlendGradient;
|
||||||
|
renderStorage->blend(mCommandEncoder, &mRenderStorageInterm, pipeline, blendMethod, opacity);
|
||||||
// restore current render pass
|
// restore current render pass
|
||||||
renderStorage->beginRenderPass(mCommandEncoder, false);
|
renderStorage->beginRenderPass(mCommandEncoder, false);
|
||||||
}
|
}
|
||||||
|
@ -186,17 +191,19 @@ bool WgRenderer::renderImage(RenderData data)
|
||||||
// use hardware blend
|
// use hardware blend
|
||||||
if (WgPipelines::isBlendMethodSupportsHW(mBlendMethod))
|
if (WgPipelines::isBlendMethodSupportsHW(mBlendMethod))
|
||||||
renderStorage->renderPicture(mContext, (WgRenderDataPicture *)data, blendType);
|
renderStorage->renderPicture(mContext, (WgRenderDataPicture *)data, blendType);
|
||||||
else { // use custom blend
|
// use custom blend
|
||||||
|
else {
|
||||||
// terminate current render pass
|
// terminate current render pass
|
||||||
renderStorage->endRenderPass();
|
renderStorage->endRenderPass();
|
||||||
// render image to render target
|
// render image to render target
|
||||||
mRenderTarget.beginRenderPass(mCommandEncoder, true);
|
mRenderStorageInterm.beginRenderPass(mCommandEncoder, true);
|
||||||
mRenderTarget.renderPicture(mContext, (WgRenderDataPicture *)data, blendType);
|
mRenderStorageInterm.renderPicture(mContext, (WgRenderDataPicture *)data, blendType);
|
||||||
mRenderTarget.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);
|
||||||
renderStorage->blend(mCommandEncoder, &mRenderTarget, blendMethod, opacity);
|
WgPipelineBlend* pipeline = &mContext.pipelines->computeBlendImage;
|
||||||
|
renderStorage->blend(mCommandEncoder, &mRenderStorageInterm, pipeline, blendMethod, opacity);
|
||||||
// restore current render pass
|
// restore current render pass
|
||||||
renderStorage->beginRenderPass(mCommandEncoder, false);
|
renderStorage->beginRenderPass(mCommandEncoder, false);
|
||||||
}
|
}
|
||||||
|
@ -293,7 +300,6 @@ bool WgRenderer::sync()
|
||||||
mContext.executeCommandEncoder(commandEncoder);
|
mContext.executeCommandEncoder(commandEncoder);
|
||||||
wgpuCommandEncoderRelease(commandEncoder);
|
wgpuCommandEncoderRelease(commandEncoder);
|
||||||
|
|
||||||
wgpuSurfacePresent(mContext.surface);
|
|
||||||
return true;
|
return true;
|
||||||
}
|
}
|
||||||
|
|
||||||
|
@ -323,10 +329,9 @@ bool WgRenderer::target(WGPUInstance instance, WGPUSurface surface, uint32_t w,
|
||||||
wgpuSurfaceConfigure(mContext.surface, &surfaceConfiguration);
|
wgpuSurfaceConfigure(mContext.surface, &surfaceConfiguration);
|
||||||
|
|
||||||
initialize();
|
initialize();
|
||||||
mRenderTarget.initialize(mContext, w, h, WG_SSAA_SAMPLES);
|
mRenderStorageInterm.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;
|
||||||
}
|
}
|
||||||
|
|
||||||
|
@ -366,11 +371,11 @@ bool WgRenderer::endComposite(TVG_UNUSED Compositor* cmp)
|
||||||
// get two last render targets
|
// get two last render targets
|
||||||
WgRenderStorage* renderStorageSrc = mRenderStorageStack.last();
|
WgRenderStorage* renderStorageSrc = mRenderStorageStack.last();
|
||||||
mRenderStorageStack.pop();
|
mRenderStorageStack.pop();
|
||||||
|
|
||||||
// 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, blendMethod, opacity);
|
mRenderStorageStack.last()->blend(mCommandEncoder, renderStorageSrc, &mContext.pipelines->computeBlendImage, blendMethod, opacity);
|
||||||
|
|
||||||
// back render targets to the pool
|
// back render targets to the pool
|
||||||
mRenderStoragePool.free(mContext, renderStorageSrc);
|
mRenderStoragePool.free(mContext, renderStorageSrc);
|
||||||
|
|
|
@ -63,13 +63,19 @@ public:
|
||||||
private:
|
private:
|
||||||
// render handles
|
// render handles
|
||||||
WGPUCommandEncoder mCommandEncoder{};
|
WGPUCommandEncoder mCommandEncoder{};
|
||||||
WgRenderStorage mRenderTarget;
|
// intermidiate buffer to render
|
||||||
|
WgRenderStorage mRenderStorageInterm;
|
||||||
|
// root render storage
|
||||||
WgRenderStorage mRenderStorageRoot;
|
WgRenderStorage mRenderStorageRoot;
|
||||||
|
// storage with data after antializing
|
||||||
WgRenderStorage mRenderStorageScreen;
|
WgRenderStorage mRenderStorageScreen;
|
||||||
|
// pool to hold render tree storages
|
||||||
WgRenderStoragePool mRenderStoragePool;
|
WgRenderStoragePool mRenderStoragePool;
|
||||||
|
// opacity, blend methods and composite methods pool
|
||||||
WgBindGroupOpacityPool mOpacityPool;
|
WgBindGroupOpacityPool mOpacityPool;
|
||||||
WgBindGroupBlendMethodPool mBlendMethodPool;
|
WgBindGroupBlendMethodPool mBlendMethodPool;
|
||||||
WgBindGroupCompositeMethodPool mCompositeMethodPool;
|
WgBindGroupCompositeMethodPool mCompositeMethodPool;
|
||||||
|
// render data shpes pool
|
||||||
WgRenderDataShapePool mRenderDataShapePool;
|
WgRenderDataShapePool mRenderDataShapePool;
|
||||||
|
|
||||||
// render tree stacks
|
// render tree stacks
|
||||||
|
@ -80,6 +86,7 @@ private:
|
||||||
WgContext mContext;
|
WgContext mContext;
|
||||||
WgPipelines mPipelines;
|
WgPipelines mPipelines;
|
||||||
Surface mTargetSurface;
|
Surface mTargetSurface;
|
||||||
|
// current blend method
|
||||||
BlendMethod mBlendMethod{};
|
BlendMethod mBlendMethod{};
|
||||||
};
|
};
|
||||||
|
|
||||||
|
|
|
@ -23,11 +23,13 @@
|
||||||
#include "tvgWgShaderSrc.h"
|
#include "tvgWgShaderSrc.h"
|
||||||
#include <string>
|
#include <string>
|
||||||
|
|
||||||
|
#define WG_SHADER_SOURCE(...) #__VA_ARGS__
|
||||||
|
|
||||||
//************************************************************************
|
//************************************************************************
|
||||||
// shader pipeline fill
|
// shader pipeline fill
|
||||||
//************************************************************************
|
//************************************************************************
|
||||||
|
|
||||||
const char* cShaderSource_PipelineFill = R"(
|
const char* cShaderSource_PipelineFill = WG_SHADER_SOURCE(
|
||||||
// vertex input
|
// vertex input
|
||||||
struct VertexInput {
|
struct VertexInput {
|
||||||
@location(0) position: vec2f
|
@location(0) position: vec2f
|
||||||
|
@ -54,13 +56,13 @@ fn vs_main(in: VertexInput) -> VertexOutput {
|
||||||
fn fs_main(in: VertexOutput) -> void {
|
fn fs_main(in: VertexOutput) -> void {
|
||||||
// nothing to draw, just stencil value
|
// nothing to draw, just stencil value
|
||||||
}
|
}
|
||||||
)";
|
);
|
||||||
|
|
||||||
//************************************************************************
|
//************************************************************************
|
||||||
// shader pipeline solid
|
// shader pipeline solid
|
||||||
//************************************************************************
|
//************************************************************************
|
||||||
|
|
||||||
const char* cShaderSource_PipelineSolid = R"(
|
const char* cShaderSource_PipelineSolid = WG_SHADER_SOURCE(
|
||||||
// vertex input
|
// vertex input
|
||||||
struct VertexInput {
|
struct VertexInput {
|
||||||
@location(0) position: vec2f
|
@location(0) position: vec2f
|
||||||
|
@ -80,10 +82,10 @@ struct VertexOutput {
|
||||||
};
|
};
|
||||||
|
|
||||||
// uniforms
|
// uniforms
|
||||||
@group(0) @binding(0) var<uniform> uViewMat : mat4x4f;
|
@group(0) @binding(0) var<uniform> uViewMat : mat4x4f;
|
||||||
@group(1) @binding(0) var<uniform> uModelMat : mat4x4f;
|
@group(1) @binding(0) var<uniform> uModelMat : mat4x4f;
|
||||||
@group(1) @binding(1) var<uniform> uBlendSettings : BlendSettings;
|
@group(1) @binding(1) var<uniform> uBlendSettings : BlendSettings;
|
||||||
@group(2) @binding(0) var<uniform> uSolidColor : vec4f;
|
@group(2) @binding(0) var<uniform> uSolidColor : vec4f;
|
||||||
|
|
||||||
@vertex
|
@vertex
|
||||||
fn vs_main(in: VertexInput) -> VertexOutput {
|
fn vs_main(in: VertexInput) -> VertexOutput {
|
||||||
|
@ -102,15 +104,15 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f {
|
||||||
color = uSolidColor;
|
color = uSolidColor;
|
||||||
|
|
||||||
let alpha: f32 = color.a * uBlendSettings.opacity;
|
let alpha: f32 = color.a * uBlendSettings.opacity;
|
||||||
return vec4f(color.rgb*alpha, alpha);
|
return vec4f(color.rgb * alpha, alpha);
|
||||||
}
|
}
|
||||||
)";
|
);
|
||||||
|
|
||||||
//************************************************************************
|
//************************************************************************
|
||||||
// shader pipeline linear
|
// shader pipeline linear
|
||||||
//************************************************************************
|
//************************************************************************
|
||||||
|
|
||||||
const char* cShaderSource_PipelineLinear = R"(
|
const char* cShaderSource_PipelineLinear = WG_SHADER_SOURCE(
|
||||||
// vertex input
|
// vertex input
|
||||||
struct VertexInput {
|
struct VertexInput {
|
||||||
@location(0) position: vec2f
|
@location(0) position: vec2f
|
||||||
|
@ -171,7 +173,7 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f {
|
||||||
let ba: vec2f = ed - st;
|
let ba: vec2f = ed - st;
|
||||||
|
|
||||||
// get interpolation factor
|
// get interpolation factor
|
||||||
var t: f32 = abs(dot(pos - st, ba) / dot(ba, ba));
|
var t: f32 = dot(pos - st, ba) / dot(ba, ba);
|
||||||
|
|
||||||
// fill spread
|
// fill spread
|
||||||
switch uLinearGradient.spread {
|
switch uLinearGradient.spread {
|
||||||
|
@ -205,15 +207,15 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f {
|
||||||
}
|
}
|
||||||
|
|
||||||
let alpha: f32 = color.a * uBlendSettings.opacity;
|
let alpha: f32 = color.a * uBlendSettings.opacity;
|
||||||
return vec4f(color.rgb*alpha, alpha);
|
return vec4f(color.rgb * alpha, alpha);
|
||||||
}
|
}
|
||||||
)";
|
);
|
||||||
|
|
||||||
//************************************************************************
|
//************************************************************************
|
||||||
// shader pipeline radial
|
// shader pipeline radial
|
||||||
//************************************************************************
|
//************************************************************************
|
||||||
|
|
||||||
const char* cShaderSource_PipelineRadial = R"(
|
const char* cShaderSource_PipelineRadial = WG_SHADER_SOURCE(
|
||||||
// vertex input
|
// vertex input
|
||||||
struct VertexInput {
|
struct VertexInput {
|
||||||
@location(0) position: vec2f
|
@location(0) position: vec2f
|
||||||
|
@ -250,7 +252,7 @@ struct VertexOutput {
|
||||||
// uniforms
|
// uniforms
|
||||||
@group(0) @binding(0) var<uniform> uViewMat : mat4x4f;
|
@group(0) @binding(0) var<uniform> uViewMat : mat4x4f;
|
||||||
@group(1) @binding(0) var<uniform> uModelMat : mat4x4f;
|
@group(1) @binding(0) var<uniform> uModelMat : mat4x4f;
|
||||||
@group(1) @binding(1) var<uniform> uBlendSettings : BlendSettings;
|
@group(1) @binding(1) var<uniform> uBlendSettings : BlendSettings;
|
||||||
@group(2) @binding(0) var<uniform> uRadialGradient : RadialGradient;
|
@group(2) @binding(0) var<uniform> uRadialGradient : RadialGradient;
|
||||||
|
|
||||||
@vertex
|
@vertex
|
||||||
|
@ -302,15 +304,15 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f {
|
||||||
}
|
}
|
||||||
|
|
||||||
let alpha: f32 = color.a * uBlendSettings.opacity;
|
let alpha: f32 = color.a * uBlendSettings.opacity;
|
||||||
return vec4f(color.rgb*alpha, alpha);
|
return vec4f(color.rgb * alpha, alpha);
|
||||||
}
|
}
|
||||||
)";
|
);
|
||||||
|
|
||||||
//************************************************************************
|
//************************************************************************
|
||||||
// cShaderSource_PipelineImage
|
// cShaderSource_PipelineImage
|
||||||
//************************************************************************
|
//************************************************************************
|
||||||
|
|
||||||
const char* cShaderSource_PipelineImage = R"(
|
const char* cShaderSource_PipelineImage = WG_SHADER_SOURCE(
|
||||||
// vertex input
|
// vertex input
|
||||||
struct VertexInput {
|
struct VertexInput {
|
||||||
@location(0) position: vec2f,
|
@location(0) position: vec2f,
|
||||||
|
@ -331,11 +333,11 @@ struct VertexOutput {
|
||||||
@location(0) texCoord: vec2f
|
@location(0) texCoord: vec2f
|
||||||
};
|
};
|
||||||
|
|
||||||
@group(0) @binding(0) var<uniform> uViewMat : mat4x4f;
|
@group(0) @binding(0) var<uniform> uViewMat : mat4x4f;
|
||||||
@group(1) @binding(0) var<uniform> uModelMat : mat4x4f;
|
@group(1) @binding(0) var<uniform> uModelMat : mat4x4f;
|
||||||
@group(1) @binding(1) var<uniform> uBlendSettings : BlendSettings;
|
@group(1) @binding(1) var<uniform> uBlendSettings : BlendSettings;
|
||||||
@group(2) @binding(0) var uSampler : sampler;
|
@group(2) @binding(0) var uSampler : sampler;
|
||||||
@group(2) @binding(1) var uTextureView : texture_2d<f32>;
|
@group(2) @binding(1) var uTextureView : texture_2d<f32>;
|
||||||
|
|
||||||
@vertex
|
@vertex
|
||||||
fn vs_main(in: VertexInput) -> VertexOutput {
|
fn vs_main(in: VertexInput) -> VertexOutput {
|
||||||
|
@ -360,25 +362,13 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f {
|
||||||
}
|
}
|
||||||
return result * uBlendSettings.opacity;
|
return result * uBlendSettings.opacity;
|
||||||
};
|
};
|
||||||
)";
|
);
|
||||||
|
|
||||||
//************************************************************************
|
//************************************************************************
|
||||||
// cShaderSource_PipelineComputeBlend
|
// cShaderSource_PipelineComputeBlend
|
||||||
//************************************************************************
|
//************************************************************************
|
||||||
|
|
||||||
// pipeline shader modules clear
|
const std::string strBlendShaderHeader = WG_SHADER_SOURCE(
|
||||||
const char* cShaderSource_PipelineComputeClear = R"(
|
|
||||||
@group(0) @binding(0) var imageDst : texture_storage_2d<rgba8unorm, read_write>;
|
|
||||||
|
|
||||||
@compute @workgroup_size(8, 8)
|
|
||||||
fn cs_main( @builtin(global_invocation_id) id: vec3u) {
|
|
||||||
textureStore(imageDst, id.xy, vec4f(0.0, 0.0, 0.0, 0.0));
|
|
||||||
}
|
|
||||||
)";
|
|
||||||
|
|
||||||
|
|
||||||
// pipeline shader modules blend
|
|
||||||
const char* cShaderSource_PipelineComputeBlend = R"(
|
|
||||||
@group(0) @binding(0) var imageSrc : texture_storage_2d<rgba8unorm, read_write>;
|
@group(0) @binding(0) var imageSrc : texture_storage_2d<rgba8unorm, read_write>;
|
||||||
@group(1) @binding(0) var imageDst : texture_storage_2d<rgba8unorm, read_write>;
|
@group(1) @binding(0) var imageDst : texture_storage_2d<rgba8unorm, read_write>;
|
||||||
@group(2) @binding(0) var<uniform> blendMethod : u32;
|
@group(2) @binding(0) var<uniform> blendMethod : u32;
|
||||||
|
@ -389,41 +379,123 @@ fn cs_main( @builtin(global_invocation_id) id: vec3u) {
|
||||||
let texSize = textureDimensions(imageSrc);
|
let texSize = textureDimensions(imageSrc);
|
||||||
if ((id.x >= texSize.x) || (id.y >= texSize.y)) { return; };
|
if ((id.x >= texSize.x) || (id.y >= texSize.y)) { return; };
|
||||||
|
|
||||||
let colorSrc = textureLoad(imageSrc, id.xy);
|
let colorSrc= textureLoad(imageSrc, id.xy);
|
||||||
|
if (colorSrc.a == 0.0) { return; }
|
||||||
let colorDst = textureLoad(imageDst, id.xy);
|
let colorDst = textureLoad(imageDst, id.xy);
|
||||||
|
|
||||||
let So: f32 = opacity;
|
var One: vec3f = vec3(1.0);
|
||||||
let Sa: f32 = colorSrc.a;
|
var So: f32 = opacity;
|
||||||
let Da: f32 = colorDst.a;
|
var Sc: vec3f = colorSrc.rgb;
|
||||||
let S: vec4f = colorSrc;
|
var Sa: f32 = colorSrc.a;
|
||||||
let D: vec4f = colorDst;
|
var Dc: vec3f = colorDst.rgb;
|
||||||
let One: vec4f = vec4(1.0);
|
var Da: f32 = colorDst.a;
|
||||||
|
var Rc: vec3f = colorDst.rgb;
|
||||||
|
var Ra: f32 = 1.0;
|
||||||
|
);
|
||||||
|
|
||||||
var color: vec4f = vec4f(0.0, 0.0, 0.0, 0.0);
|
const std::string strBlendShaderPreConditionsGradient = WG_SHADER_SOURCE(
|
||||||
|
Sc = Sc + Dc.rgb * (1.0 - Sa);
|
||||||
|
Sa = Sa + Da * (1.0 - Sa);
|
||||||
|
);
|
||||||
|
|
||||||
|
const std::string strBlendShaderPreConditionsImage = WG_SHADER_SOURCE(
|
||||||
|
Sc = Sc * So;
|
||||||
|
Sa = Sa * So;
|
||||||
|
);
|
||||||
|
|
||||||
|
const std::string strBlendShaderBlendMethod = WG_SHADER_SOURCE(
|
||||||
switch blendMethod {
|
switch blendMethod {
|
||||||
/* Normal */ case 0u: { color = (S * So) + (1.0 - Sa * So) * D; }
|
/* Normal */ case 0u: {
|
||||||
/* Add */ case 1u: { color = (S + D); }
|
Rc = Sc + Dc * (1.0 - Sa);
|
||||||
/* Screen */ case 2u: { color = (S + D) - (S * D); }
|
Ra = Sa + Da * (1.0 - Sa);
|
||||||
/* Multiply */ case 3u: { color = (S * D); }
|
}
|
||||||
/* Overlay */ case 4u: { color = (Sa * Da) - 2 * (Da - S) * (Sa - D); }
|
/* Add */ case 1u: { Rc = Sc + Dc; }
|
||||||
/* Difference */ case 5u: { color = abs(S - D); }
|
/* Screen */ case 2u: { Rc = Sc + Dc - Sc * Dc; }
|
||||||
/* Exclusion */ case 6u: { color = S + D - (2 * S * D); }
|
/* Multiply */ case 3u: { Rc = Sc * Dc; }
|
||||||
/* SrcOver */ case 7u: { color = S; }
|
/* Overlay */ case 4u: {
|
||||||
/* Darken */ case 8u: { color = min(S, D); }
|
Rc.r = select(1.0 - min(1.0, 2 * (1 - Sc.r) * (1 - Dc.r)), min(1.0, 2 * Sc.r * Dc.r), (Dc.r < 0.5));
|
||||||
/* Lighten */ case 9u: { color = max(S, D); }
|
Rc.g = select(1.0 - min(1.0, 2 * (1 - Sc.g) * (1 - Dc.g)), min(1.0, 2 * Sc.g * Dc.g), (Dc.g < 0.5));
|
||||||
/* ColorDodge */ case 10u: { color = D / (One - S); }
|
Rc.b = select(1.0 - min(1.0, 2 * (1 - Sc.b) * (1 - Dc.b)), min(1.0, 2 * Sc.b * Dc.b), (Dc.b < 0.5));
|
||||||
/* ColorBurn */ case 11u: { color = One - (One - D) / S; }
|
}
|
||||||
/* HardLight */ case 12u: { color = (Sa * Da) - 2.0 * (Da - S) * (Sa - D); }
|
/* Difference */ case 5u: { Rc = abs(Dc - Sc); }
|
||||||
/* SoftLight */ case 13u: { color = (One - 2 * S) * (D * D) + (2 * S * D); }
|
/* Exclusion */ case 6u: { Rc = min(One, Sc + Dc - min(One, 2 * Sc * Dc)); }
|
||||||
default: { color = (S * So) + (1.0 - Sa * So) * D; }
|
/* SrcOver */ case 7u: { Rc = Sc; Ra = Sa; }
|
||||||
|
/* Darken */ case 8u: { Rc = min(Sc, Dc); }
|
||||||
|
/* Lighten */ case 9u: { Rc = max(Sc, Dc); }
|
||||||
|
/* ColorDodge */ case 10u: {
|
||||||
|
Rc.r = select(Dc.r, (Dc.r * 255.0 / (255.0 - Sc.r * 255.0))/255.0, (1.0 - Sc.r > 0.0));
|
||||||
|
Rc.g = select(Dc.g, (Dc.g * 255.0 / (255.0 - Sc.g * 255.0))/255.0, (1.0 - Sc.g > 0.0));
|
||||||
|
Rc.b = select(Dc.b, (Dc.b * 255.0 / (255.0 - Sc.b * 255.0))/255.0, (1.0 - Sc.b > 0.0));
|
||||||
|
}
|
||||||
|
/* ColorBurn */ case 11u: {
|
||||||
|
Rc.r = select(1.0 - Dc.r, (255.0 - (255.0 - Dc.r * 255.0) / (Sc.r * 255.0)) / 255.0, (Sc.r > 0.0));
|
||||||
|
Rc.g = select(1.0 - Dc.g, (255.0 - (255.0 - Dc.g * 255.0) / (Sc.g * 255.0)) / 255.0, (Sc.g > 0.0));
|
||||||
|
Rc.b = select(1.0 - Dc.b, (255.0 - (255.0 - Dc.b * 255.0) / (Sc.b * 255.0)) / 255.0, (Sc.b > 0.0));
|
||||||
|
}
|
||||||
|
/* HardLight */ case 12u: {
|
||||||
|
Rc.r = select(1.0 - min(1.0, 2 * (1 - Sc.r) * (1 - Dc.r)), min(1.0, 2 * Sc.r * Dc.r), (Sc.r < 0.5));
|
||||||
|
Rc.g = select(1.0 - min(1.0, 2 * (1 - Sc.g) * (1 - Dc.g)), min(1.0, 2 * Sc.g * Dc.g), (Sc.g < 0.5));
|
||||||
|
Rc.b = select(1.0 - min(1.0, 2 * (1 - Sc.b) * (1 - Dc.b)), min(1.0, 2 * Sc.b * Dc.b), (Sc.b < 0.5));
|
||||||
|
}
|
||||||
|
/* SoftLight */ case 13u: { Rc = min(One, (One - 2 * Sc) * Dc * Dc + 2.0 * Sc * Dc); }
|
||||||
|
default: {
|
||||||
|
Rc = Sc + Dc * (1.0 - Sa);
|
||||||
|
Ra = Sa + Da * (1.0 - Sa);
|
||||||
|
}
|
||||||
}
|
}
|
||||||
|
);
|
||||||
|
|
||||||
textureStore(imageDst, id.xy, color);
|
const std::string strBlendShaderPostConditionsGradient = WG_SHADER_SOURCE(
|
||||||
|
// nothing
|
||||||
|
);
|
||||||
|
|
||||||
|
const std::string strBlendShaderPostConditionsImage = WG_SHADER_SOURCE(
|
||||||
|
Rc = select(mix(Dc, Rc, Sa), Rc, blendMethod == 0);
|
||||||
|
Ra = select(mix(Da, Ra, Sa), Ra, blendMethod == 0);
|
||||||
|
);
|
||||||
|
|
||||||
|
const std::string strBlendShaderFooter = WG_SHADER_SOURCE(
|
||||||
|
textureStore(imageDst, id.xy, vec4(Rc, Ra));
|
||||||
}
|
}
|
||||||
)";
|
);
|
||||||
|
|
||||||
|
// pipeline shader modules blend solid
|
||||||
|
const std::string strComputeBlendSolid =
|
||||||
|
strBlendShaderHeader +
|
||||||
|
strBlendShaderBlendMethod +
|
||||||
|
strBlendShaderFooter;
|
||||||
|
const char* cShaderSource_PipelineComputeBlendSolid = strComputeBlendSolid.c_str();
|
||||||
|
|
||||||
|
// pipeline shader modules blend gradient
|
||||||
|
const std::string strComputeBlendGradient =
|
||||||
|
strBlendShaderHeader +
|
||||||
|
strBlendShaderPreConditionsGradient +
|
||||||
|
strBlendShaderBlendMethod +
|
||||||
|
strBlendShaderPostConditionsGradient +
|
||||||
|
strBlendShaderFooter;
|
||||||
|
const char* cShaderSource_PipelineComputeBlendGradient = strComputeBlendGradient.c_str();
|
||||||
|
|
||||||
|
// pipeline shader modules blend image
|
||||||
|
const std::string strComputeBlendImage =
|
||||||
|
strBlendShaderHeader +
|
||||||
|
strBlendShaderPreConditionsImage +
|
||||||
|
strBlendShaderBlendMethod +
|
||||||
|
strBlendShaderPostConditionsImage +
|
||||||
|
strBlendShaderFooter;
|
||||||
|
const char* cShaderSource_PipelineComputeBlendImage = strComputeBlendImage.c_str();
|
||||||
|
|
||||||
|
// pipeline shader modules clear
|
||||||
|
const char* cShaderSource_PipelineComputeClear = WG_SHADER_SOURCE(
|
||||||
|
@group(0) @binding(0) var imageDst : texture_storage_2d<rgba8unorm, read_write>;
|
||||||
|
|
||||||
|
@compute @workgroup_size(8, 8)
|
||||||
|
fn cs_main( @builtin(global_invocation_id) id: vec3u) {
|
||||||
|
textureStore(imageDst, id.xy, vec4f(0.0, 0.0, 0.0, 0.0));
|
||||||
|
}
|
||||||
|
);
|
||||||
|
|
||||||
// pipeline shader modules compose
|
// pipeline shader modules compose
|
||||||
const char* cShaderSource_PipelineComputeCompose = R"(
|
const char* cShaderSource_PipelineComputeCompose = WG_SHADER_SOURCE(
|
||||||
@group(0) @binding(0) var imageSrc : texture_storage_2d<rgba8unorm, read_write>;
|
@group(0) @binding(0) var imageSrc : texture_storage_2d<rgba8unorm, read_write>;
|
||||||
@group(1) @binding(0) var imageMsk : texture_storage_2d<rgba8unorm, read_write>;
|
@group(1) @binding(0) var imageMsk : texture_storage_2d<rgba8unorm, read_write>;
|
||||||
@group(2) @binding(0) var<uniform> composeMethod : u32;
|
@group(2) @binding(0) var<uniform> composeMethod : u32;
|
||||||
|
@ -455,10 +527,10 @@ fn cs_main( @builtin(global_invocation_id) id: vec3u) {
|
||||||
|
|
||||||
textureStore(imageSrc, id.xy, vec4f(color, alpha * opacity));
|
textureStore(imageSrc, id.xy, vec4f(color, alpha * opacity));
|
||||||
}
|
}
|
||||||
)";
|
);
|
||||||
|
|
||||||
// pipeline shader modules compose blend
|
// pipeline shader modules compose blend
|
||||||
const char* cShaderSource_PipelineComputeComposeBlend = R"(
|
const char* cShaderSource_PipelineComputeComposeBlend = 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>;
|
||||||
|
@ -521,10 +593,10 @@ fn cs_main( @builtin(global_invocation_id) id: vec3u) {
|
||||||
|
|
||||||
textureStore(imageDst, id.xy, color);
|
textureStore(imageDst, id.xy, color);
|
||||||
}
|
}
|
||||||
)";
|
);
|
||||||
|
|
||||||
// pipeline shader modules anti-aliasing
|
// pipeline shader modules anti-aliasing
|
||||||
const char* cShaderSource_PipelineComputeAntiAlias = R"(
|
const char* cShaderSource_PipelineComputeAntiAlias = WG_SHADER_SOURCE(
|
||||||
@group(0) @binding(0) var imageSrc : texture_storage_2d<rgba8unorm, read_write>;
|
@group(0) @binding(0) var imageSrc : texture_storage_2d<rgba8unorm, read_write>;
|
||||||
@group(1) @binding(0) var imageDst : texture_storage_2d<bgra8unorm, write>;
|
@group(1) @binding(0) var imageDst : texture_storage_2d<bgra8unorm, write>;
|
||||||
|
|
||||||
|
@ -544,4 +616,4 @@ fn cs_main( @builtin(global_invocation_id) id: vec3u) {
|
||||||
|
|
||||||
textureStore(imageDst, id.xy, color / f32(samples * samples));
|
textureStore(imageDst, id.xy, color / f32(samples * samples));
|
||||||
}
|
}
|
||||||
)";
|
);
|
||||||
|
|
|
@ -42,7 +42,9 @@ extern const char* cShaderSource_PipelineImage;
|
||||||
|
|
||||||
// pipeline shader modules clear, compose and blend
|
// pipeline shader modules clear, compose and blend
|
||||||
extern const char* cShaderSource_PipelineComputeClear;
|
extern const char* cShaderSource_PipelineComputeClear;
|
||||||
extern const char* cShaderSource_PipelineComputeBlend;
|
extern const char* cShaderSource_PipelineComputeBlendSolid;
|
||||||
|
extern const char* cShaderSource_PipelineComputeBlendGradient;
|
||||||
|
extern const char* cShaderSource_PipelineComputeBlendImage;
|
||||||
extern const char* cShaderSource_PipelineComputeCompose;
|
extern const char* cShaderSource_PipelineComputeCompose;
|
||||||
extern const char* cShaderSource_PipelineComputeComposeBlend;
|
extern const char* cShaderSource_PipelineComputeComposeBlend;
|
||||||
extern const char* cShaderSource_PipelineComputeAntiAlias;
|
extern const char* cShaderSource_PipelineComputeAntiAlias;
|
||||||
|
|
Loading…
Add table
Reference in a new issue