mirror of
https://github.com/thorvg/thorvg.git
synced 2025-06-08 13:43:43 +00:00
wg_engine: aabb based masking optimization
* used fragment shaders for mask applience instead of compute shaders * less render targets swithing * shape aabb based on transformed shape bbox
This commit is contained in:
Sergii Liebodkin
Hermet Park
parent
8598348ec0
commit
0787e46635
10 changed files with 267 additions and 219 deletions
|
|
@ -47,11 +47,17 @@ void WgCompositor::initialize(WgContext& context, uint32_t width, uint32_t heigh
|
|||
storageClipPath.initialize(context, width, height);
|
||||
storageInterm.initialize(context, width, height);
|
||||
storageDstCopy.initialize(context, width, height);
|
||||
// composition and blend geometries
|
||||
WgGeometryData geometryData;
|
||||
geometryData.appendBlitBox();
|
||||
meshData.update(context, &geometryData);
|
||||
}
|
||||
|
||||
|
||||
void WgCompositor::release(WgContext& context)
|
||||
{
|
||||
// composition and blend geometries
|
||||
meshData.release(context);
|
||||
// release intermediate render storages
|
||||
storageInterm.release(context);
|
||||
storageDstCopy.release(context);
|
||||
|
|
@ -252,6 +258,7 @@ void WgCompositor::blendImage(WgContext& context, WgRenderDataPicture* renderDat
|
|||
};
|
||||
|
||||
|
||||
// TODO: use direct mask applience
|
||||
void WgCompositor::composeShape(WgContext& context, WgRenderDataShape* renderData, WgRenderStorage* mask, CompositeMethod composeMethod)
|
||||
{
|
||||
assert(mask);
|
||||
|
|
@ -268,16 +275,17 @@ void WgCompositor::composeShape(WgContext& context, WgRenderDataShape* renderDat
|
|||
beginRenderPass(commandEncoder, &storageInterm, true);
|
||||
drawShape(context, renderData, WgPipelineBlendType::Custom);
|
||||
endRenderPass();
|
||||
// run compisition
|
||||
compose(commandEncoder, &storageInterm, mask, target, composeMethod);
|
||||
// restore current render pass
|
||||
beginRenderPass(commandEncoder, target, false);
|
||||
RenderRegion rect { 0, 0,(int32_t)width, (int32_t)height };
|
||||
composeRegion(context, &storageInterm, mask, composeMethod, rect);
|
||||
}
|
||||
|
||||
|
||||
void WgCompositor::composeStrokes(WgContext& context, WgRenderDataShape* renderData, WgRenderStorage* msk, CompositeMethod composeMethod)
|
||||
// TODO: use direct mask applience
|
||||
void WgCompositor::composeStrokes(WgContext& context, WgRenderDataShape* renderData, WgRenderStorage* mask, CompositeMethod composeMethod)
|
||||
{
|
||||
assert(msk);
|
||||
assert(mask);
|
||||
assert(renderData);
|
||||
assert(commandEncoder);
|
||||
assert(currentTarget);
|
||||
|
|
@ -291,16 +299,17 @@ void WgCompositor::composeStrokes(WgContext& context, WgRenderDataShape* renderD
|
|||
beginRenderPass(commandEncoder, &storageInterm, true);
|
||||
drawStrokes(context, renderData, WgPipelineBlendType::Custom);
|
||||
endRenderPass();
|
||||
// run compisition
|
||||
compose(commandEncoder, &storageInterm, msk, target, composeMethod);
|
||||
// restore current render pass
|
||||
beginRenderPass(commandEncoder, target, false);
|
||||
RenderRegion rect { 0, 0, (int32_t)width, (int32_t)height };
|
||||
composeRegion(context, &storageInterm, mask, composeMethod, rect);
|
||||
}
|
||||
|
||||
|
||||
void WgCompositor::composeImage(WgContext& context, WgRenderDataPicture* renderData, WgRenderStorage* msk, CompositeMethod composeMethod)
|
||||
// TODO: use direct mask applience
|
||||
void WgCompositor::composeImage(WgContext& context, WgRenderDataPicture* renderData, WgRenderStorage* mask, CompositeMethod composeMethod)
|
||||
{
|
||||
assert(msk);
|
||||
assert(mask);
|
||||
assert(renderData);
|
||||
assert(commandEncoder);
|
||||
assert(currentTarget);
|
||||
|
|
@ -311,10 +320,38 @@ void WgCompositor::composeImage(WgContext& context, WgRenderDataPicture* renderD
|
|||
beginRenderPass(commandEncoder, &storageInterm, true);
|
||||
drawImage(context, renderData, WgPipelineBlendType::Custom);
|
||||
endRenderPass();
|
||||
// run compisition
|
||||
compose(commandEncoder, &storageInterm, msk, target, composeMethod);
|
||||
// restore current render pass
|
||||
beginRenderPass(commandEncoder, target, false);
|
||||
RenderRegion rect { 0, 0, (int32_t)width, (int32_t)height };
|
||||
composeRegion(context, &storageInterm, mask, composeMethod, rect);
|
||||
}
|
||||
|
||||
|
||||
void WgCompositor::composeScene(WgContext& context, WgRenderStorage* src, WgRenderStorage* mask, WgCompose* cmp)
|
||||
{
|
||||
assert(cmp);
|
||||
assert(src);
|
||||
assert(mask);
|
||||
assert(renderPassEncoder);
|
||||
// cut viewport to screen dimensions
|
||||
int32_t xmin = std::max(0, std::min((int32_t)width, cmp->aabb.x));
|
||||
int32_t ymin = std::max(0, std::min((int32_t)height, cmp->aabb.y));
|
||||
int32_t xmax = std::max(0, std::min((int32_t)width, cmp->aabb.x + cmp->aabb.w));
|
||||
int32_t ymax = std::max(0, std::min((int32_t)height, cmp->aabb.y + cmp->aabb.h));
|
||||
if ((xmin >= xmax) || (ymin >= ymax)) return;
|
||||
RenderRegion rect { xmin, ymin, xmax - xmin, ymax - ymin };
|
||||
composeRegion(context, src, mask, cmp->method, rect);
|
||||
}
|
||||
|
||||
|
||||
void WgCompositor::composeRegion(WgContext& context, WgRenderStorage* src, WgRenderStorage* mask, CompositeMethod composeMethod, RenderRegion& rect)
|
||||
{
|
||||
wgpuRenderPassEncoderSetScissorRect(renderPassEncoder, rect.x, rect.y, rect.w, rect.h);
|
||||
wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 0);
|
||||
wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 0, src->bindGroupTexure, 0, nullptr);
|
||||
wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 1, mask->bindGroupTexure, 0, nullptr);
|
||||
wgpuRenderPassEncoderSetPipeline(renderPassEncoder, pipelines->sceneComp[(uint32_t)composeMethod]);
|
||||
meshData.drawImage(context, renderPassEncoder);
|
||||
}
|
||||
|
||||
|
||||
|
|
@ -483,27 +520,3 @@ void WgCompositor::blend(WGPUCommandEncoder encoder, WgRenderStorage* src, WgRen
|
|||
wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, (width + 7) / 8, (height + 7) / 8, 1);
|
||||
wgpuComputePassEncoderEnd(computePassEncoder);
|
||||
}
|
||||
|
||||
|
||||
void WgCompositor::compose(WGPUCommandEncoder encoder, WgRenderStorage* src, WgRenderStorage* mask, WgRenderStorage* dst, CompositeMethod composeMethod)
|
||||
{
|
||||
assert(src);
|
||||
assert(mask);
|
||||
assert(dst);
|
||||
assert(!renderPassEncoder);
|
||||
// copy dst storage to temporary read only storage
|
||||
const WGPUImageCopyTexture texSrc { .texture = dst->texture };
|
||||
const WGPUImageCopyTexture texDst { .texture = storageDstCopy.texture };
|
||||
const WGPUExtent3D copySize { .width = width, .height = height, .depthOrArrayLayers = 1 };
|
||||
wgpuCommandEncoderCopyTextureToTexture(encoder, &texSrc, &texDst, ©Size);
|
||||
// execute compose shader
|
||||
const WGPUComputePassDescriptor computePassDescriptor{};
|
||||
WGPUComputePassEncoder computePassEncoder = wgpuCommandEncoderBeginComputePass(encoder, &computePassDescriptor);
|
||||
wgpuComputePassEncoderSetBindGroup(computePassEncoder, 0, src->bindGroupRead, 0, nullptr);
|
||||
wgpuComputePassEncoderSetBindGroup(computePassEncoder, 1, mask->bindGroupRead, 0, nullptr);
|
||||
wgpuComputePassEncoderSetBindGroup(computePassEncoder, 2, storageDstCopy.bindGroupRead, 0, nullptr);
|
||||
wgpuComputePassEncoderSetBindGroup(computePassEncoder, 3, dst->bindGroupWrite, 0, nullptr);
|
||||
wgpuComputePassEncoderSetPipeline(computePassEncoder, pipelines->compose[(size_t)composeMethod]);
|
||||
wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, (width + 7) / 8, (height + 7) / 8, 1);
|
||||
wgpuComputePassEncoderEnd(computePassEncoder);
|
||||
}
|
||||
|
|
|
|||
|
|
@ -27,7 +27,8 @@
|
|||
#include "tvgWgRenderData.h"
|
||||
|
||||
struct WgCompose: public Compositor {
|
||||
BlendMethod blend;
|
||||
BlendMethod blend{};
|
||||
RenderRegion aabb{};
|
||||
};
|
||||
|
||||
class WgCompositor {
|
||||
|
|
@ -51,8 +52,12 @@ private:
|
|||
WgRenderStorage storageInterm;
|
||||
WgRenderStorage storageClipPath;
|
||||
WgRenderStorage storageDstCopy;
|
||||
// composition and blend geometries
|
||||
WgMeshData meshData;
|
||||
|
||||
static WgPipelineBlendType blendMethodToBlendType(BlendMethod blendMethod);
|
||||
|
||||
void composeRegion(WgContext& context, WgRenderStorage* src, WgRenderStorage* mask, CompositeMethod composeMethod, RenderRegion& rect);
|
||||
public:
|
||||
// render target dimensions
|
||||
uint32_t width{};
|
||||
|
|
@ -76,6 +81,7 @@ public:
|
|||
void composeShape(WgContext& context, WgRenderDataShape* renderData, WgRenderStorage* mask, CompositeMethod composeMethod);
|
||||
void composeStrokes(WgContext& context, WgRenderDataShape* renderData, WgRenderStorage* mask, CompositeMethod composeMethod);
|
||||
void composeImage(WgContext& context, WgRenderDataPicture* renderData, WgRenderStorage* mask, CompositeMethod composeMethod);
|
||||
void composeScene(WgContext& context, WgRenderStorage* src, WgRenderStorage* mask, WgCompose* cmp);
|
||||
|
||||
void drawClipPath(WgContext& context, WgRenderDataShape* renderData);
|
||||
void drawShape(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType);
|
||||
|
|
@ -84,7 +90,6 @@ public:
|
|||
|
||||
void mergeMasks(WGPUCommandEncoder encoder, WgRenderStorage* mask0, WgRenderStorage* mask1);
|
||||
void blend(WGPUCommandEncoder encoder, WgRenderStorage* src, WgRenderStorage* dst, uint8_t opacity, BlendMethod blendMethod, WgRenderRasterType rasterType);
|
||||
void compose(WGPUCommandEncoder encoder, WgRenderStorage* src, WgRenderStorage* mask, WgRenderStorage* dst, CompositeMethod composeMethod);
|
||||
};
|
||||
|
||||
#endif // _TVG_WG_COMPOSITOR_H_
|
||||
|
|
|
|||
|
|
@ -59,7 +59,7 @@ public:
|
|||
inline void normalize() { float rlen = 1.0f / length(); x *= rlen; y *= rlen; }
|
||||
inline WgPoint normal() const { float rlen = 1.0f / length(); return { x * rlen, y * rlen }; }
|
||||
inline WgPoint lerp(const WgPoint& p, float t) const { return { x + (p.x - x) * t, y + (p.y - y) * t }; };
|
||||
inline WgPoint trans(const Matrix& m) const { return { x * m.e11 + y * m.e12, x * m.e21 + y * m.e22 }; };
|
||||
inline WgPoint trans(const Matrix& m) const { return { x * m.e11 + y * m.e12 + m.e13, x * m.e21 + y * m.e22 + m.e23 }; };
|
||||
};
|
||||
|
||||
struct WgMath
|
||||
|
|
|
|||
|
|
@ -43,7 +43,8 @@ WGPUPipelineLayout WgPipelines::createPipelineLayout(WGPUDevice device, const WG
|
|||
|
||||
WGPURenderPipeline WgPipelines::createRenderPipeline(
|
||||
WGPUDevice device, const char* pipelineLabel,
|
||||
const WGPUShaderModule shaderModule, const WGPUPipelineLayout pipelineLayout,
|
||||
const WGPUShaderModule shaderModule, const char* vsEntryPoint, const char* fsEntryPoint,
|
||||
const WGPUPipelineLayout pipelineLayout,
|
||||
const WGPUVertexBufferLayout *vertexBufferLayouts, const uint32_t vertexBufferLayoutsCount,
|
||||
const WGPUColorWriteMaskFlags writeMask,
|
||||
const WGPUCompareFunction stencilFunctionFrnt, const WGPUStencilOperation stencilOperationFrnt,
|
||||
|
|
@ -58,8 +59,8 @@ WGPURenderPipeline WgPipelines::createRenderPipeline(
|
|||
.stencilBack = { .compare = stencilFunctionBack, .failOp = stencilOperationBack, .depthFailOp = stencilOperationBack, .passOp = stencilOperationBack },
|
||||
.stencilReadMask = 0xFFFFFFFF, .stencilWriteMask = 0xFFFFFFFF
|
||||
};
|
||||
const WGPUVertexState vertexState { .module = shaderModule, .entryPoint = "vs_main", .bufferCount = vertexBufferLayoutsCount, .buffers = vertexBufferLayouts };
|
||||
const WGPUFragmentState fragmentState { .module = shaderModule, .entryPoint = "fs_main", .targetCount = 1, .targets = colorTargetStates };
|
||||
const WGPUVertexState vertexState { .module = shaderModule, .entryPoint = vsEntryPoint, .bufferCount = vertexBufferLayoutsCount, .buffers = vertexBufferLayouts };
|
||||
const WGPUFragmentState fragmentState { .module = shaderModule, .entryPoint = fsEntryPoint, .targetCount = 1, .targets = colorTargetStates };
|
||||
const WGPURenderPipelineDescriptor renderPipelineDesc {
|
||||
.label = pipelineLabel,
|
||||
.layout = pipelineLayout,
|
||||
|
|
@ -182,6 +183,10 @@ void WgPipelines::initialize(WgContext& context)
|
|||
layouts.layoutBuffer2Un,
|
||||
layouts.layoutTexSampled
|
||||
};
|
||||
const WGPUBindGroupLayout bindGroupLayoutsSceneComp[] = {
|
||||
layouts.layoutTexSampled,
|
||||
layouts.layoutTexSampled
|
||||
};
|
||||
const WGPUBindGroupLayout bindGroupLayoutsMergeMasks[] = {
|
||||
layouts.layoutTexStrorage1RO,
|
||||
layouts.layoutTexStrorage1RO,
|
||||
|
|
@ -193,12 +198,6 @@ void WgPipelines::initialize(WgContext& context)
|
|||
layouts.layoutTexStrorage1WO,
|
||||
layouts.layoutBuffer1Un
|
||||
};
|
||||
const WGPUBindGroupLayout bindGroupLayoutsCompose[] = {
|
||||
layouts.layoutTexStrorage1RO,
|
||||
layouts.layoutTexStrorage1RO,
|
||||
layouts.layoutTexStrorage1RO,
|
||||
layouts.layoutTexStrorage1WO
|
||||
};
|
||||
const WGPUBindGroupLayout bindGroupLayoutsCopy[] = {
|
||||
layouts.layoutTexStrorage1RO,
|
||||
layouts.layoutTexScreen1WO
|
||||
|
|
@ -209,8 +208,8 @@ void WgPipelines::initialize(WgContext& context)
|
|||
layoutSolid = createPipelineLayout(context.device, bindGroupLayoutsSolid, 3);
|
||||
layoutGradient = createPipelineLayout(context.device, bindGroupLayoutsGradient, 3);
|
||||
layoutImage = createPipelineLayout(context.device, bindGroupLayoutsImage, 3);
|
||||
layoutSceneComp = createPipelineLayout(context.device, bindGroupLayoutsSceneComp, 2);
|
||||
layoutBlend = createPipelineLayout(context.device, bindGroupLayoutsBlend, 4);
|
||||
layoutCompose = createPipelineLayout(context.device, bindGroupLayoutsCompose, 4);
|
||||
layoutMergeMasks = createPipelineLayout(context.device, bindGroupLayoutsMergeMasks, 3);
|
||||
layoutCopy = createPipelineLayout(context.device, bindGroupLayoutsCopy, 2);
|
||||
|
||||
|
|
@ -220,6 +219,7 @@ void WgPipelines::initialize(WgContext& context)
|
|||
shaderRadial = createShaderModule(context.device, "The shader radial", cShaderSrc_Radial);
|
||||
shaderLinear = createShaderModule(context.device, "The shader linear", cShaderSrc_Linear);
|
||||
shaderImage = createShaderModule(context.device, "The shader image", cShaderSrc_Image);
|
||||
shaderSceneComp = createShaderModule(context.device, "The shader scene composition", cShaderSrc_Scene_Comp);
|
||||
// computes shader modules
|
||||
shaderMergeMasks = createShaderModule(context.device, "The shader merge mask", cShaderSrc_MergeMasks);
|
||||
shaderCopy = createShaderModule(context.device, "The shader copy", cShaderSrc_Copy);
|
||||
|
|
@ -227,7 +227,7 @@ void WgPipelines::initialize(WgContext& context)
|
|||
// render pipeline winding
|
||||
winding = createRenderPipeline(
|
||||
context.device, "The render pipeline winding",
|
||||
shaderStencil, layoutStencil,
|
||||
shaderStencil, "vs_main", "fs_main", layoutStencil,
|
||||
vertexBufferLayoutsShape, 1,
|
||||
WGPUColorWriteMask_None,
|
||||
WGPUCompareFunction_Always, WGPUStencilOperation_IncrementWrap,
|
||||
|
|
@ -236,7 +236,7 @@ void WgPipelines::initialize(WgContext& context)
|
|||
// render pipeline even-odd
|
||||
evenodd = createRenderPipeline(
|
||||
context.device, "The render pipeline even-odd",
|
||||
shaderStencil, layoutStencil,
|
||||
shaderStencil, "vs_main", "fs_main", layoutStencil,
|
||||
vertexBufferLayoutsShape, 1,
|
||||
WGPUColorWriteMask_None,
|
||||
WGPUCompareFunction_Always, WGPUStencilOperation_Invert,
|
||||
|
|
@ -245,7 +245,7 @@ void WgPipelines::initialize(WgContext& context)
|
|||
// render pipeline direct
|
||||
direct = createRenderPipeline(
|
||||
context.device, "The render pipeline direct",
|
||||
shaderStencil, layoutStencil,
|
||||
shaderStencil,"vs_main", "fs_main", layoutStencil,
|
||||
vertexBufferLayoutsShape, 1,
|
||||
WGPUColorWriteMask_None,
|
||||
WGPUCompareFunction_Always, WGPUStencilOperation_Replace,
|
||||
|
|
@ -254,7 +254,7 @@ void WgPipelines::initialize(WgContext& context)
|
|||
// render pipeline clip path
|
||||
clipPath = createRenderPipeline(
|
||||
context.device, "The render pipeline clip path",
|
||||
shaderStencil, layoutStencil,
|
||||
shaderStencil, "vs_main", "fs_main", layoutStencil,
|
||||
vertexBufferLayoutsShape, 1,
|
||||
WGPUColorWriteMask_All,
|
||||
WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero,
|
||||
|
|
@ -265,7 +265,7 @@ void WgPipelines::initialize(WgContext& context)
|
|||
for (uint32_t i = 0; i < 3; i++) {
|
||||
solid[i] = createRenderPipeline(
|
||||
context.device, "The render pipeline solid",
|
||||
shaderSolid, layoutSolid,
|
||||
shaderSolid, "vs_main", "fs_main", layoutSolid,
|
||||
vertexBufferLayoutsShape, 1,
|
||||
WGPUColorWriteMask_All,
|
||||
WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero,
|
||||
|
|
@ -277,7 +277,7 @@ void WgPipelines::initialize(WgContext& context)
|
|||
for (uint32_t i = 0; i < 3; i++) {
|
||||
radial[i] = createRenderPipeline(
|
||||
context.device, "The render pipeline radial",
|
||||
shaderRadial, layoutGradient,
|
||||
shaderRadial, "vs_main", "fs_main", layoutGradient,
|
||||
vertexBufferLayoutsShape, 1,
|
||||
WGPUColorWriteMask_All,
|
||||
WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero,
|
||||
|
|
@ -289,7 +289,7 @@ void WgPipelines::initialize(WgContext& context)
|
|||
for (uint32_t i = 0; i < 3; i++) {
|
||||
linear[i] = createRenderPipeline(
|
||||
context.device, "The render pipeline linear",
|
||||
shaderLinear, layoutGradient,
|
||||
shaderLinear, "vs_main", "fs_main", layoutGradient,
|
||||
vertexBufferLayoutsShape, 1,
|
||||
WGPUColorWriteMask_All,
|
||||
WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero,
|
||||
|
|
@ -301,7 +301,7 @@ void WgPipelines::initialize(WgContext& context)
|
|||
for (uint32_t i = 0; i < 3; i++) {
|
||||
image[i] = createRenderPipeline(
|
||||
context.device, "The render pipeline image",
|
||||
shaderImage, layoutImage,
|
||||
shaderImage, "vs_main", "fs_main", layoutImage,
|
||||
vertexBufferLayoutsImage, 2,
|
||||
WGPUColorWriteMask_All,
|
||||
WGPUCompareFunction_Always, WGPUStencilOperation_Zero,
|
||||
|
|
@ -309,6 +309,55 @@ void WgPipelines::initialize(WgContext& context)
|
|||
primitiveState, multisampleState, blendStates[i]);
|
||||
}
|
||||
|
||||
// compute shader names
|
||||
const char* shaderComposeNames[] {
|
||||
"fs_main_None",
|
||||
"fs_main_ClipPath",
|
||||
"fs_main_AlphaMask",
|
||||
"fs_main_InvAlphaMask",
|
||||
"fs_main_LumaMask",
|
||||
"fs_main_InvLumaMask",
|
||||
"fs_main_AddMask",
|
||||
"fs_main_SubtractMask",
|
||||
"fs_main_IntersectMask",
|
||||
"fs_main_DifferenceMask",
|
||||
"fs_main_LightenMask",
|
||||
"fs_main_DarkenMask"
|
||||
};
|
||||
|
||||
// compose shader blend states
|
||||
const WGPUBlendState composeBlends[] {
|
||||
blendStateNrm, // None
|
||||
blendStateNrm, // ClipPath
|
||||
blendStateNrm, // AlphaMask
|
||||
blendStateNrm, // InvAlphaMask
|
||||
blendStateNrm, // LumaMask
|
||||
blendStateNrm, // InvLumaMask
|
||||
blendStateSrc, // AddMask
|
||||
blendStateSrc, // SubtractMask
|
||||
blendStateSrc, // IntersectMask
|
||||
blendStateSrc, // DifferenceMask
|
||||
blendStateSrc, // LightenMask
|
||||
blendStateSrc // DarkenMask
|
||||
};
|
||||
|
||||
// render pipeline scene composition
|
||||
size_t shaderComposeNamesCnt = sizeof(shaderComposeNames) / sizeof(shaderComposeNames[0]);
|
||||
size_t composeBlendspCnt = sizeof(composeBlends) / sizeof(composeBlends[0]);
|
||||
size_t sceneCompCnt = sizeof(sceneComp) / sizeof(sceneComp[0]);
|
||||
assert(shaderComposeNamesCnt == composeBlendspCnt);
|
||||
assert(composeBlendspCnt == sceneCompCnt);
|
||||
for (uint32_t i = 0; i < sceneCompCnt; i++) {
|
||||
sceneComp[i] = createRenderPipeline(
|
||||
context.device, "The render pipeline scene composition",
|
||||
shaderSceneComp, "vs_main", shaderComposeNames[i], layoutSceneComp,
|
||||
vertexBufferLayoutsImage, 2,
|
||||
WGPUColorWriteMask_All,
|
||||
WGPUCompareFunction_Always, WGPUStencilOperation_Zero,
|
||||
WGPUCompareFunction_Always, WGPUStencilOperation_Zero,
|
||||
primitiveState, multisampleState, composeBlends[i]);
|
||||
}
|
||||
|
||||
// compute pipelines
|
||||
mergeMasks = createComputePipeline(context.device, "The pipeline merge masks", shaderMergeMasks, "cs_main", layoutMergeMasks);
|
||||
copy = createComputePipeline(context.device, "The pipeline copy", shaderCopy, "cs_main", layoutCopy);
|
||||
|
|
@ -346,37 +395,14 @@ void WgPipelines::initialize(WgContext& context)
|
|||
blendGradient[i] = createComputePipeline(context.device, "The pipeline blend gradient", shaderBlendGradient, shaderBlendNames[i], layoutBlend);
|
||||
blendImage[i] = createComputePipeline(context.device, "The pipeline blend image", shaderBlendImage, shaderBlendNames[i], layoutBlend);
|
||||
}
|
||||
|
||||
// compute shader compose names
|
||||
const char* shaderComposeNames[] {
|
||||
"cs_main_None",
|
||||
"cs_main_ClipPath",
|
||||
"cs_main_AlphaMask",
|
||||
"cs_main_InvAlphaMask",
|
||||
"cs_main_LumaMask",
|
||||
"cs_main_InvLumaMask",
|
||||
"cs_main_AddMask",
|
||||
"cs_main_SubtractMask",
|
||||
"cs_main_IntersectMask",
|
||||
"cs_main_DifferenceMask",
|
||||
"cs_main_LightenMask",
|
||||
"cs_main_DarkenMask"
|
||||
};
|
||||
|
||||
// create compose shaders
|
||||
shaderCompose = createShaderModule(context.device, "The shader compose", cShaderSrc_Compose);
|
||||
|
||||
// create compose pipelines
|
||||
const size_t shaderComposeNamesCnt = sizeof(shaderComposeNames) / sizeof(shaderComposeNames[0]);
|
||||
const size_t pipesComposeCnt = sizeof(compose) / sizeof(compose[0]);
|
||||
assert(shaderComposeNamesCnt == pipesComposeCnt);
|
||||
for (uint32_t i = 0; i < pipesComposeCnt; i++)
|
||||
compose[i] = createComputePipeline(context.device, "The pipeline compose", shaderCompose, shaderComposeNames[i], layoutCompose);
|
||||
}
|
||||
|
||||
void WgPipelines::releaseGraphicHandles(WgContext& context)
|
||||
{
|
||||
releaseRenderPipeline(clipPath);
|
||||
size_t pipesSceneCompCnt = sizeof(sceneComp) / sizeof(sceneComp[0]);
|
||||
for (uint32_t i = 0; i < pipesSceneCompCnt; i++)
|
||||
releaseRenderPipeline(sceneComp[i]);
|
||||
for (uint32_t i = 0; i < 3; i++) {
|
||||
releaseRenderPipeline(image[i]);
|
||||
releaseRenderPipeline(linear[i]);
|
||||
|
|
@ -386,10 +412,12 @@ void WgPipelines::releaseGraphicHandles(WgContext& context)
|
|||
releaseRenderPipeline(direct);
|
||||
releaseRenderPipeline(evenodd);
|
||||
releaseRenderPipeline(winding);
|
||||
releasePipelineLayout(layoutSceneComp);
|
||||
releasePipelineLayout(layoutImage);
|
||||
releasePipelineLayout(layoutGradient);
|
||||
releasePipelineLayout(layoutSolid);
|
||||
releasePipelineLayout(layoutStencil);
|
||||
releaseShaderModule(shaderSceneComp);
|
||||
releaseShaderModule(shaderImage);
|
||||
releaseShaderModule(shaderLinear);
|
||||
releaseShaderModule(shaderRadial);
|
||||
|
|
@ -400,9 +428,6 @@ void WgPipelines::releaseGraphicHandles(WgContext& context)
|
|||
|
||||
void WgPipelines::releaseComputeHandles(WgContext& context)
|
||||
{
|
||||
size_t pipesComposeCnt = sizeof(compose) / sizeof(compose[0]);
|
||||
for (uint32_t i = 0; i < pipesComposeCnt; i++)
|
||||
releaseComputePipeline(compose[i]);
|
||||
const size_t pipesBlendCnt = sizeof(blendSolid)/sizeof(blendSolid[0]);
|
||||
for (uint32_t i = 0; i < pipesBlendCnt; i++) {
|
||||
releaseComputePipeline(blendImage[i]);
|
||||
|
|
@ -411,11 +436,9 @@ void WgPipelines::releaseComputeHandles(WgContext& context)
|
|||
}
|
||||
releaseComputePipeline(copy);
|
||||
releaseComputePipeline(mergeMasks);
|
||||
releasePipelineLayout(layoutCompose);
|
||||
releasePipelineLayout(layoutBlend);
|
||||
releasePipelineLayout(layoutMergeMasks);
|
||||
releasePipelineLayout(layoutCopy);
|
||||
releaseShaderModule(shaderCompose);
|
||||
releaseShaderModule(shaderBlendImage);
|
||||
releaseShaderModule(shaderBlendGradient);
|
||||
releaseShaderModule(shaderBlendSolid);
|
||||
|
|
|
|||
|
|
@ -35,12 +35,12 @@ private:
|
|||
WGPUShaderModule shaderRadial{};
|
||||
WGPUShaderModule shaderLinear{};
|
||||
WGPUShaderModule shaderImage{};
|
||||
WGPUShaderModule shaderSceneComp{};
|
||||
// compute pipeline shaders
|
||||
WGPUShaderModule shaderMergeMasks;
|
||||
WGPUShaderModule shaderBlendSolid;
|
||||
WGPUShaderModule shaderBlendGradient;
|
||||
WGPUShaderModule shaderBlendImage;
|
||||
WGPUShaderModule shaderCompose;
|
||||
WGPUShaderModule shaderCopy;
|
||||
private:
|
||||
// graphics pipeline layouts
|
||||
|
|
@ -48,10 +48,10 @@ private:
|
|||
WGPUPipelineLayout layoutSolid{};
|
||||
WGPUPipelineLayout layoutGradient{};
|
||||
WGPUPipelineLayout layoutImage{};
|
||||
WGPUPipelineLayout layoutSceneComp{};
|
||||
// compute pipeline layouts
|
||||
WGPUPipelineLayout layoutMergeMasks{};
|
||||
WGPUPipelineLayout layoutBlend{};
|
||||
WGPUPipelineLayout layoutCompose{};
|
||||
WGPUPipelineLayout layoutCopy{};
|
||||
public:
|
||||
// graphics pipeline
|
||||
|
|
@ -63,13 +63,13 @@ public:
|
|||
WGPURenderPipeline radial[3]{};
|
||||
WGPURenderPipeline linear[3]{};
|
||||
WGPURenderPipeline image[3]{};
|
||||
WGPURenderPipeline sceneComp[12];
|
||||
WGPURenderPipeline clipPath{};
|
||||
// compute pipeline
|
||||
WGPUComputePipeline mergeMasks;
|
||||
WGPUComputePipeline blendSolid[14];
|
||||
WGPUComputePipeline blendGradient[14];
|
||||
WGPUComputePipeline blendImage[14];
|
||||
WGPUComputePipeline compose[12];
|
||||
WGPUComputePipeline copy;
|
||||
private:
|
||||
void releaseGraphicHandles(WgContext& context);
|
||||
|
|
@ -79,7 +79,8 @@ private:
|
|||
WGPUPipelineLayout createPipelineLayout(WGPUDevice device, const WGPUBindGroupLayout* bindGroupLayouts, const uint32_t bindGroupLayoutsCount);
|
||||
WGPURenderPipeline createRenderPipeline(
|
||||
WGPUDevice device, const char* pipelineLabel,
|
||||
const WGPUShaderModule shaderModule, const WGPUPipelineLayout pipelineLayout,
|
||||
const WGPUShaderModule shaderModule, const char* vsEntryPoint, const char* fsEntryPoint,
|
||||
const WGPUPipelineLayout pipelineLayout,
|
||||
const WGPUVertexBufferLayout *vertexBufferLayouts, const uint32_t vertexBufferLayoutsCount,
|
||||
const WGPUColorWriteMaskFlags writeMask,
|
||||
const WGPUCompareFunction stencilFunctionFrnt, const WGPUStencilOperation stencilOperationFrnt,
|
||||
|
|
|
|||
|
|
@ -306,6 +306,18 @@ void WgRenderDataShape::updateBBox(WgPoint pmin, WgPoint pmax)
|
|||
}
|
||||
|
||||
|
||||
void WgRenderDataShape::updateAABB(const Matrix& rt) {
|
||||
WgPoint p0 = WgPoint(pMin.x, pMin.y).trans(rt);
|
||||
WgPoint p1 = WgPoint(pMax.x, pMin.y).trans(rt);
|
||||
WgPoint p2 = WgPoint(pMin.x, pMax.y).trans(rt);
|
||||
WgPoint p3 = WgPoint(pMax.x, pMax.y).trans(rt);
|
||||
aabb.x = std::min({p0.x, p1.x, p2.x, p3.x});
|
||||
aabb.y = std::min({p0.y, p1.y, p2.y, p3.y});
|
||||
aabb.w = std::max({p0.x, p1.x, p2.x, p3.x}) - aabb.x;
|
||||
aabb.h = std::max({p0.y, p1.y, p2.y, p3.y}) - aabb.y;
|
||||
}
|
||||
|
||||
|
||||
void WgRenderDataShape::updateMeshes(WgContext &context, const RenderShape &rshape, const Matrix& rt)
|
||||
{
|
||||
releaseMeshes(context);
|
||||
|
|
@ -368,6 +380,7 @@ void WgRenderDataShape::updateMeshes(WgContext &context, const RenderShape &rsha
|
|||
for (uint32_t i = 0; i < polylines.count; i++)
|
||||
delete polylines[i];
|
||||
// update shapes bbox
|
||||
updateAABB(rt);
|
||||
meshDataBBox.update(context, pMin, pMax);
|
||||
}
|
||||
|
||||
|
|
@ -443,6 +456,7 @@ void WgRenderDataShape::updateStrokes(WgContext& context, const WgPolyline* poly
|
|||
geometryData.positions.getBBox(pmin, pmax);
|
||||
meshGroupStrokes.append(context, &geometryData);
|
||||
meshGroupStrokesBBox.append(context, pmin, pmax);
|
||||
updateBBox(pmin, pmax);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
|
|
|||
|
|
@ -97,6 +97,7 @@ struct WgRenderDataPaint
|
|||
WGPUBuffer bufferBlendSettings{};
|
||||
WGPUBindGroup bindGroupPaint{};
|
||||
RenderRegion viewport{};
|
||||
RenderRegion aabb{};
|
||||
float opacity{};
|
||||
Array<WgRenderDataPaint*> clips;
|
||||
|
||||
|
|
@ -123,6 +124,7 @@ struct WgRenderDataShape: public WgRenderDataPaint
|
|||
FillRule fillRule{};
|
||||
|
||||
void updateBBox(WgPoint pmin, WgPoint pmax);
|
||||
void updateAABB(const Matrix& rt);
|
||||
void updateMeshes(WgContext& context, const RenderShape& rshape, const Matrix& rt);
|
||||
void updateShapes(WgContext& context, const WgPolyline* polyline);
|
||||
void updateStrokesList(WgContext& context, Array<WgPolyline*> polylines, const RenderStroke* rstroke, float totalLen, float trimBegin, float trimEnd);
|
||||
|
|
|
|||
|
|
@ -199,8 +199,12 @@ void WgRenderer::dispose(RenderData data) {
|
|||
}
|
||||
|
||||
|
||||
RenderRegion WgRenderer::region(TVG_UNUSED RenderData data)
|
||||
RenderRegion WgRenderer::region(RenderData data)
|
||||
{
|
||||
auto renderData = (WgRenderDataPaint*)data;
|
||||
if (renderData->type() == Type::Shape) {
|
||||
return renderData->aabb;
|
||||
}
|
||||
return { 0, 0, (int32_t)mTargetSurface.w, (int32_t)mTargetSurface.h };
|
||||
}
|
||||
|
||||
|
|
@ -316,6 +320,7 @@ bool WgRenderer::target(WGPUInstance instance, WGPUSurface surface, uint32_t w,
|
|||
Compositor* WgRenderer::target(TVG_UNUSED const RenderRegion& region, TVG_UNUSED ColorSpace cs)
|
||||
{
|
||||
mCompositorStack.push(new WgCompose);
|
||||
mCompositorStack.last()->aabb = region;
|
||||
return mCompositorStack.last();
|
||||
}
|
||||
|
||||
|
|
@ -354,6 +359,8 @@ bool WgRenderer::endComposite(Compositor* cmp)
|
|||
mCompositor.blend(mCommandEncoder, src, dst, comp->opacity, comp->blend, WgRenderRasterType::Image);
|
||||
// back render targets to the pool
|
||||
mRenderStoragePool.free(mContext, src);
|
||||
// begin previous render pass
|
||||
mCompositor.beginRenderPass(mCommandEncoder, mRenderStorageStack.last(), false);
|
||||
} else { // finish composition
|
||||
// get source, mask and destination render storages
|
||||
WgRenderStorage* src = mRenderStorageStack.last();
|
||||
|
|
@ -361,14 +368,14 @@ bool WgRenderer::endComposite(Compositor* cmp)
|
|||
WgRenderStorage* msk = mRenderStorageStack.last();
|
||||
mRenderStorageStack.pop();
|
||||
WgRenderStorage* dst = mRenderStorageStack.last();
|
||||
// begin previous render pass
|
||||
mCompositor.beginRenderPass(mCommandEncoder, dst, false);
|
||||
// apply composition
|
||||
mCompositor.compose(mCommandEncoder, src, msk, dst,comp->method);
|
||||
mCompositor.composeScene(mContext, src, msk, comp);
|
||||
// back render targets to the pool
|
||||
mRenderStoragePool.free(mContext, src);
|
||||
mRenderStoragePool.free(mContext, msk);
|
||||
}
|
||||
// begin previous render pass
|
||||
mCompositor.beginRenderPass(mCommandEncoder, mRenderStorageStack.last(), false);
|
||||
|
||||
// delete current compositor settings
|
||||
delete mCompositorStack.last();
|
||||
|
|
|
|||
|
|
@ -184,6 +184,113 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f {
|
|||
};
|
||||
)";
|
||||
|
||||
//************************************************************************
|
||||
// graphics shader source: scene compose
|
||||
//************************************************************************
|
||||
|
||||
const char* cShaderSrc_Scene_Comp = R"(
|
||||
struct VertexInput { @location(0) position: vec2f, @location(1) texCoord: vec2f };
|
||||
struct VertexOutput { @builtin(position) position: vec4f, @location(0) texCoord: vec2f };
|
||||
|
||||
@group(0) @binding(0) var uSamplerSrc : sampler;
|
||||
@group(0) @binding(1) var uTextureSrc : texture_2d<f32>;
|
||||
@group(1) @binding(0) var uSamplerMsk : sampler;
|
||||
@group(1) @binding(1) var uTextureMsk : texture_2d<f32>;
|
||||
|
||||
@vertex
|
||||
fn vs_main(in: VertexInput) -> VertexOutput {
|
||||
var out: VertexOutput;
|
||||
out.position = vec4f(in.position.xy, 0.0, 1.0);
|
||||
out.texCoord = in.texCoord;
|
||||
return out;
|
||||
}
|
||||
|
||||
@fragment
|
||||
fn fs_main_None(in: VertexOutput) -> @location(0) vec4f {
|
||||
let src: vec4f = textureSample(uTextureSrc, uSamplerSrc, in.texCoord.xy);
|
||||
return vec4f(src);
|
||||
};
|
||||
|
||||
@fragment
|
||||
fn fs_main_ClipPath(in: VertexOutput) -> @location(0) vec4f {
|
||||
let src: vec4f = textureSample(uTextureSrc, uSamplerSrc, in.texCoord.xy);
|
||||
let msk: vec4f = textureSample(uTextureMsk, uSamplerMsk, in.texCoord.xy);
|
||||
return vec4f(src * msk.a);
|
||||
};
|
||||
|
||||
@fragment
|
||||
fn fs_main_AlphaMask(in: VertexOutput) -> @location(0) vec4f {
|
||||
let src: vec4f = textureSample(uTextureSrc, uSamplerSrc, in.texCoord.xy);
|
||||
let msk: vec4f = textureSample(uTextureMsk, uSamplerMsk, in.texCoord.xy);
|
||||
return vec4f(src * msk.a);
|
||||
};
|
||||
|
||||
@fragment
|
||||
fn fs_main_InvAlphaMask(in: VertexOutput) -> @location(0) vec4f {
|
||||
let src: vec4f = textureSample(uTextureSrc, uSamplerSrc, in.texCoord.xy);
|
||||
let msk: vec4f = textureSample(uTextureMsk, uSamplerMsk, in.texCoord.xy);
|
||||
return vec4f(src * (1.0 - msk.a));
|
||||
};
|
||||
|
||||
@fragment
|
||||
fn fs_main_LumaMask(in: VertexOutput) -> @location(0) vec4f {
|
||||
let src: vec4f = textureSample(uTextureSrc, uSamplerSrc, in.texCoord.xy);
|
||||
let msk: vec4f = textureSample(uTextureMsk, uSamplerMsk, in.texCoord.xy);
|
||||
let luma: f32 = dot(msk.rgb, vec3f(0.2125, 0.7154, 0.0721));
|
||||
return vec4f(src * luma);
|
||||
};
|
||||
|
||||
@fragment
|
||||
fn fs_main_InvLumaMask(in: VertexOutput) -> @location(0) vec4f {
|
||||
let src: vec4f = textureSample(uTextureSrc, uSamplerSrc, in.texCoord.xy);
|
||||
let msk: vec4f = textureSample(uTextureMsk, uSamplerMsk, in.texCoord.xy);
|
||||
let luma: f32 = dot(msk.rgb, vec3f(0.2125, 0.7154, 0.0721));
|
||||
return vec4f(src * (1.0 - luma));
|
||||
};
|
||||
|
||||
@fragment
|
||||
fn fs_main_AddMask(in: VertexOutput) -> @location(0) vec4f {
|
||||
let src: vec4f = textureSample(uTextureSrc, uSamplerSrc, in.texCoord.xy);
|
||||
let msk: vec4f = textureSample(uTextureMsk, uSamplerMsk, in.texCoord.xy);
|
||||
return vec4f(src.rgb, src.a + msk.a * (1.0 - src.a));
|
||||
};
|
||||
|
||||
@fragment
|
||||
fn fs_main_SubtractMask(in: VertexOutput) -> @location(0) vec4f {
|
||||
let src: vec4f = textureSample(uTextureSrc, uSamplerSrc, in.texCoord.xy);
|
||||
let msk: vec4f = textureSample(uTextureMsk, uSamplerMsk, in.texCoord.xy);
|
||||
return vec4f(src.rgb, src.a * (1.0 - msk.a));
|
||||
};
|
||||
|
||||
@fragment
|
||||
fn fs_main_IntersectMask(in: VertexOutput) -> @location(0) vec4f {
|
||||
let src: vec4f = textureSample(uTextureSrc, uSamplerSrc, in.texCoord.xy);
|
||||
let msk: vec4f = textureSample(uTextureMsk, uSamplerMsk, in.texCoord.xy);
|
||||
return vec4f(src.rgb, src.a * msk.a);
|
||||
};
|
||||
|
||||
@fragment
|
||||
fn fs_main_DifferenceMask(in: VertexOutput) -> @location(0) vec4f {
|
||||
let src: vec4f = textureSample(uTextureSrc, uSamplerSrc, in.texCoord.xy);
|
||||
let msk: vec4f = textureSample(uTextureMsk, uSamplerMsk, in.texCoord.xy);
|
||||
return vec4f(src.rgb, src.a * (1.0 - msk.a) + msk.a * (1.0 - src.a));
|
||||
};
|
||||
|
||||
@fragment
|
||||
fn fs_main_LightenMask(in: VertexOutput) -> @location(0) vec4f {
|
||||
let src: vec4f = textureSample(uTextureSrc, uSamplerSrc, in.texCoord.xy);
|
||||
let msk: vec4f = textureSample(uTextureMsk, uSamplerMsk, in.texCoord.xy);
|
||||
return vec4f(src.rgb, max(src.a, msk.a));
|
||||
};
|
||||
|
||||
@fragment
|
||||
fn fs_main_DarkenMask(in: VertexOutput) -> @location(0) vec4f {
|
||||
let src: vec4f = textureSample(uTextureSrc, uSamplerSrc, in.texCoord.xy);
|
||||
let msk: vec4f = textureSample(uTextureMsk, uSamplerMsk, in.texCoord.xy);
|
||||
return vec4f(src.rgb, min(src.a, msk.a));
|
||||
};
|
||||
)";
|
||||
|
||||
//************************************************************************
|
||||
// compute shader source: merge clip path masks
|
||||
//************************************************************************
|
||||
|
|
@ -412,130 +519,6 @@ fn cs_main_SoftLight(@builtin(global_invocation_id) id: vec3u) {
|
|||
};
|
||||
)";
|
||||
|
||||
//************************************************************************
|
||||
// compute shader source: compose
|
||||
//************************************************************************
|
||||
|
||||
const char* cShaderSrc_Compose = R"(
|
||||
@group(0) @binding(0) var imageSrc : texture_storage_2d<rgba8unorm, read>;
|
||||
@group(1) @binding(0) var imageMsk : texture_storage_2d<rgba8unorm, read>;
|
||||
@group(2) @binding(0) var imageDst : texture_storage_2d<rgba8unorm, read>;
|
||||
@group(3) @binding(0) var imageTgt : texture_storage_2d<rgba8unorm, write>;
|
||||
|
||||
struct FragData { Sc: vec3f, Sa: f32, Mc: vec3f, Ma: f32, Dc: vec3f, Da: f32 };
|
||||
fn getFragData(id: vec2u) -> FragData {
|
||||
let colorSrc = textureLoad(imageSrc, id.xy);
|
||||
let colorMsk = textureLoad(imageMsk, id.xy);
|
||||
let colorDst = textureLoad(imageDst, id.xy);
|
||||
var data: FragData;
|
||||
data.Sc = colorSrc.rgb;
|
||||
data.Sa = colorSrc.a;
|
||||
data.Mc = colorMsk.rgb;
|
||||
data.Ma = colorMsk.a;
|
||||
data.Dc = colorDst.rgb;
|
||||
data.Da = colorDst.a;
|
||||
return data;
|
||||
};
|
||||
|
||||
@compute @workgroup_size(8, 8)
|
||||
fn cs_main_None(@builtin(global_invocation_id) id: vec3u) {
|
||||
let d: FragData = getFragData(id.xy);
|
||||
let Rc = d.Dc;
|
||||
let Ra = d.Da;
|
||||
textureStore(imageTgt, id.xy, vec4f(Rc, Ra));
|
||||
};
|
||||
|
||||
@compute @workgroup_size(8, 8)
|
||||
fn cs_main_ClipPath(@builtin(global_invocation_id) id: vec3u) {
|
||||
let d: FragData = getFragData(id.xy);
|
||||
let Rc = d.Sc * d.Ma + d.Dc * (1.0 - d.Sa * d.Ma);
|
||||
let Ra = d.Sa * d.Ma + d.Da * (1.0 - d.Sa * d.Ma);
|
||||
textureStore(imageTgt, id.xy, vec4f(Rc, Ra));
|
||||
};
|
||||
|
||||
@compute @workgroup_size(8, 8)
|
||||
fn cs_main_AlphaMask(@builtin(global_invocation_id) id: vec3u) {
|
||||
let d: FragData = getFragData(id.xy);
|
||||
let Rc = d.Sc * d.Ma + d.Dc * (1.0 - d.Sa * d.Ma);
|
||||
let Ra = d.Sa * d.Ma + d.Da * (1.0 - d.Sa * d.Ma);
|
||||
textureStore(imageTgt, id.xy, vec4f(Rc, Ra));
|
||||
};
|
||||
|
||||
@compute @workgroup_size(8, 8)
|
||||
fn cs_main_InvAlphaMask(@builtin(global_invocation_id) id: vec3u) {
|
||||
let d: FragData = getFragData(id.xy);
|
||||
let Rc = d.Sc * (1.0 - d.Ma) + d.Dc * (1.0 - d.Sa * (1.0 - d.Ma));
|
||||
let Ra = d.Sa * (1.0 - d.Ma) + d.Da * (1.0 - d.Sa * (1.0 - d.Ma));
|
||||
textureStore(imageTgt, id.xy, vec4f(Rc, Ra));
|
||||
};
|
||||
|
||||
@compute @workgroup_size(8, 8)
|
||||
fn cs_main_LumaMask(@builtin(global_invocation_id) id: vec3u) {
|
||||
let d: FragData = getFragData(id.xy);
|
||||
let luma: f32 = dot(d.Mc, vec3f(0.2125, 0.7154, 0.0721));
|
||||
let Rc = d.Sc * luma + d.Dc * (1.0 - d.Sa * luma);
|
||||
let Ra = d.Sa * luma + d.Da * (1.0 - d.Sa * luma);
|
||||
textureStore(imageTgt, id.xy, vec4f(Rc, Ra));
|
||||
};
|
||||
|
||||
@compute @workgroup_size(8, 8)
|
||||
fn cs_main_InvLumaMask(@builtin(global_invocation_id) id: vec3u) {
|
||||
let d: FragData = getFragData(id.xy);
|
||||
let luma: f32 = dot(d.Mc, vec3f(0.2125, 0.7154, 0.0721));
|
||||
let Rc = d.Sc * (1.0 - luma) + d.Dc * (1.0 - d.Sa * (1.0 - luma));
|
||||
let Ra = d.Sa * (1.0 - luma) + d.Da * (1.0 - d.Sa * (1.0 - luma));
|
||||
textureStore(imageTgt, id.xy, vec4f(Rc, Ra));
|
||||
};
|
||||
|
||||
@compute @workgroup_size(8, 8)
|
||||
fn cs_main_AddMask(@builtin(global_invocation_id) id: vec3u) {
|
||||
let d: FragData = getFragData(id.xy);
|
||||
let Rc = d.Sc;
|
||||
let Ra = d.Sa + d.Ma * (1.0 - d.Sa);
|
||||
textureStore(imageTgt, id.xy, vec4f(Rc, Ra));
|
||||
};
|
||||
|
||||
@compute @workgroup_size(8, 8)
|
||||
fn cs_main_SubtractMask(@builtin(global_invocation_id) id: vec3u) {
|
||||
let d: FragData = getFragData(id.xy);
|
||||
let Rc = d.Sc;
|
||||
let Ra = d.Sa * (1.0 - d.Ma);
|
||||
textureStore(imageTgt, id.xy, vec4f(Rc, Ra));
|
||||
};
|
||||
|
||||
@compute @workgroup_size(8, 8)
|
||||
fn cs_main_IntersectMask(@builtin(global_invocation_id) id: vec3u) {
|
||||
let d: FragData = getFragData(id.xy);
|
||||
let Rc = d.Sc;
|
||||
let Ra = d.Sa * d.Ma;
|
||||
textureStore(imageTgt, id.xy, vec4f(Rc, Ra));
|
||||
};
|
||||
|
||||
@compute @workgroup_size(8, 8)
|
||||
fn cs_main_DifferenceMask(@builtin(global_invocation_id) id: vec3u) {
|
||||
let d: FragData = getFragData(id.xy);
|
||||
let Rc = d.Sc;
|
||||
let Ra = d.Sa * (1.0 - d.Ma) + d.Ma * (1.0 - d.Sa);
|
||||
textureStore(imageTgt, id.xy, vec4f(Rc, Ra));
|
||||
};
|
||||
|
||||
@compute @workgroup_size(8, 8)
|
||||
fn cs_main_LightenMask(@builtin(global_invocation_id) id: vec3u) {
|
||||
let d: FragData = getFragData(id.xy);
|
||||
let Rc = d.Sc;
|
||||
let Ra = max(d.Sa, d.Ma);
|
||||
textureStore(imageTgt, id.xy, vec4f(Rc, Ra));
|
||||
};
|
||||
|
||||
@compute @workgroup_size(8, 8)
|
||||
fn cs_main_DarkenMask(@builtin(global_invocation_id) id: vec3u) {
|
||||
let d: FragData = getFragData(id.xy);
|
||||
let Rc = d.Sc;
|
||||
let Ra = min(d.Sa, d.Ma);
|
||||
textureStore(imageTgt, id.xy, vec4f(Rc, Ra));
|
||||
};
|
||||
)";
|
||||
|
||||
//************************************************************************
|
||||
// compute shader source: copy
|
||||
//************************************************************************
|
||||
|
|
|
|||
|
|
@ -29,6 +29,7 @@ extern const char* cShaderSrc_Solid;
|
|||
extern const char* cShaderSrc_Linear;
|
||||
extern const char* cShaderSrc_Radial;
|
||||
extern const char* cShaderSrc_Image;
|
||||
extern const char* cShaderSrc_Scene_Comp;
|
||||
|
||||
// compute shader sources: blend, compose and merge path
|
||||
extern const char* cShaderSrc_MergeMasks;
|
||||
|
|
@ -36,7 +37,6 @@ extern const char* cShaderSrc_BlendHeader_Solid;
|
|||
extern const char* cShaderSrc_BlendHeader_Gradient;
|
||||
extern const char* cShaderSrc_BlendHeader_Image;
|
||||
extern const char* cShaderSrc_Blend_Funcs;
|
||||
extern const char* cShaderSrc_Compose;
|
||||
extern const char* cShaderSrc_Copy;
|
||||
|
||||
#endif // _TVG_WG_SHEDER_SRC_H_
|
||||
|
|
|
|||
Loading…
Add table
Reference in a new issue