git_clone/thorvg
1
0
Fork 0
mirror of https://github.com/thorvg/thorvg.git synced 2025-06-14 12:04:29 +00:00
Code Issues Projects Releases Packages Wiki Activity Actions 5

wg_engine: introduced blending

Browse source 
[issues 1479: blending](#1479)

Supported blend settings:

    Normal
    Add
    Screen
    Multiply
    Overlay
    Difference
    Exclusion
    SrcOver
    Darken
    Lighten
    ColorDodge
    ColorBurn
    HardLight
    SoftLight
This commit is contained in:
Sergii Liebodkin 2024-01-28 23:10:54 +02:00 • committed by Hermet Park
parent 2aa2ac7e5c
commit 9e0cc0298c
14 changed files with 850 additions and 669 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

252
src/renderer/wg_engine/tvgWgBindGroups.cpp
View file

@ -32,10 +32,12 @@ WGPUBindGroupLayout WgBindGroupLinearGradient::layout = nullptr;
WGPUBindGroupLayout WgBindGroupRadialGradient::layout = nullptr; WGPUBindGroupLayout WgBindGroupRadialGradient::layout = nullptr;
WGPUBindGroupLayout WgBindGroupPicture::layout = nullptr; WGPUBindGroupLayout WgBindGroupPicture::layout = nullptr;
// composition and blending properties gropus // composition and blending properties gropus
WGPUBindGroupLayout WgBindGroupOpacity::layout = nullptr;
WGPUBindGroupLayout WgBindGroupTexture::layout = nullptr; WGPUBindGroupLayout WgBindGroupTexture::layout = nullptr;
WGPUBindGroupLayout WgBindGroupStorageTexture::layout = nullptr; WGPUBindGroupLayout WgBindGroupTextureStorage::layout = nullptr;
WGPUBindGroupLayout WgBindGroupTextureSampled::layout = nullptr; WGPUBindGroupLayout WgBindGroupTextureSampled::layout = nullptr;
WGPUBindGroupLayout WgBindGroupOpacity::layout = nullptr;
WGPUBindGroupLayout WgBindGroupBlendMethod::layout = nullptr;
WGPUBindGroupLayout WgBindGroupCompositeMethod::layout = nullptr;
WGPUBindGroupLayout WgBindGroupCanvas::getLayout(WGPUDevice device) WGPUBindGroupLayout WgBindGroupCanvas::getLayout(WGPUDevice device)
@ -264,50 +266,6 @@ void WgBindGroupPicture::release()
} }
WGPUBindGroupLayout WgBindGroupOpacity::getLayout(WGPUDevice device)
{
if (layout) return layout;
const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] {
makeBindGroupLayoutEntryBuffer(0)
};
layout = createBindGroupLayout(device, bindGroupLayoutEntries, 1);
assert(layout);
return layout;
}
void WgBindGroupOpacity::releaseLayout()
{
releaseBindGroupLayout(layout);
}
void WgBindGroupOpacity::initialize(WGPUDevice device, WGPUQueue queue, uint32_t uOpacity)
{
release();
float opacity = uOpacity / 255.0f;
uBufferOpacity = createBuffer(device, queue, &opacity, sizeof(float));
const WGPUBindGroupEntry bindGroupEntries[] {
makeBindGroupEntryBuffer(0, uBufferOpacity)
};
mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 1);
assert(mBindGroup);
}
void WgBindGroupOpacity::update(WGPUDevice device, WGPUQueue queue, uint32_t uOpacity) {
float opacity = uOpacity / 255.0f;
wgpuQueueWriteBuffer(queue, uBufferOpacity, 0, &opacity, sizeof(float));
}
void WgBindGroupOpacity::release()
{
releaseBuffer(uBufferOpacity);
releaseBindGroup(mBindGroup);
}
WGPUBindGroupLayout WgBindGroupTexture::getLayout(WGPUDevice device) WGPUBindGroupLayout WgBindGroupTexture::getLayout(WGPUDevice device)
{ {
if (layout) return layout; if (layout) return layout;
@ -343,11 +301,11 @@ void WgBindGroupTexture::release()
} }
WGPUBindGroupLayout WgBindGroupStorageTexture::getLayout(WGPUDevice device) WGPUBindGroupLayout WgBindGroupTextureStorage::getLayout(WGPUDevice device)
{ {
if (layout) return layout; if (layout) return layout;
const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] {
makeBindGroupLayoutEntryStorageTexture(0) makeBindGroupLayoutEntryStorageTexture(0, WGPUStorageTextureAccess_ReadWrite)
}; };
layout = createBindGroupLayout(device, bindGroupLayoutEntries, 1); layout = createBindGroupLayout(device, bindGroupLayoutEntries, 1);
assert(layout); assert(layout);
@ -355,13 +313,13 @@ WGPUBindGroupLayout WgBindGroupStorageTexture::getLayout(WGPUDevice device)
} }
void WgBindGroupStorageTexture::releaseLayout() void WgBindGroupTextureStorage::releaseLayout()
{ {
releaseBindGroupLayout(layout); releaseBindGroupLayout(layout);
} }
void WgBindGroupStorageTexture::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexture) void WgBindGroupTextureStorage::initialize(WGPUDevice device, WGPUQueue queue, WGPUTextureView uTexture)
{ {
release(); release();
const WGPUBindGroupEntry bindGroupEntries[] { const WGPUBindGroupEntry bindGroupEntries[] {
@ -372,7 +330,7 @@ void WgBindGroupStorageTexture::initialize(WGPUDevice device, WGPUQueue queue, W
} }
void WgBindGroupStorageTexture::release() void WgBindGroupTextureStorage::release()
{ {
releaseBindGroup(mBindGroup); releaseBindGroup(mBindGroup);
} }
@ -414,13 +372,131 @@ void WgBindGroupTextureSampled::release()
releaseBindGroup(mBindGroup); releaseBindGroup(mBindGroup);
} }
WGPUBindGroupLayout WgBindGroupOpacity::getLayout(WGPUDevice device)
{
if (layout) return layout;
const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] {
makeBindGroupLayoutEntryBuffer(0)
};
layout = createBindGroupLayout(device, bindGroupLayoutEntries, 1);
assert(layout);
return layout;
}
void WgBindGroupOpacity::releaseLayout()
{
releaseBindGroupLayout(layout);
}
void WgBindGroupOpacity::initialize(WGPUDevice device, WGPUQueue queue, uint32_t uOpacity)
{
release();
float opacity = uOpacity / 255.0f;
uBufferOpacity = createBuffer(device, queue, &opacity, sizeof(float));
const WGPUBindGroupEntry bindGroupEntries[] {
makeBindGroupEntryBuffer(0, uBufferOpacity)
};
mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 1);
assert(mBindGroup);
}
void WgBindGroupOpacity::release()
{
releaseBuffer(uBufferOpacity);
releaseBindGroup(mBindGroup);
}
WGPUBindGroupLayout WgBindGroupBlendMethod::getLayout(WGPUDevice device)
{
if (layout) return layout;
const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] {
makeBindGroupLayoutEntryBuffer(0)
};
layout = createBindGroupLayout(device, bindGroupLayoutEntries, 1);
assert(layout);
return layout;
}
void WgBindGroupBlendMethod::releaseLayout()
{
releaseBindGroupLayout(layout);
}
void WgBindGroupBlendMethod::initialize(WGPUDevice device, WGPUQueue queue, BlendMethod uBlendMethod)
{
release();
uint32_t blendMethod = (uint32_t)uBlendMethod;
uBufferBlendMethod = createBuffer(device, queue, &blendMethod, sizeof(blendMethod));
const WGPUBindGroupEntry bindGroupEntries[] {
makeBindGroupEntryBuffer(0, uBufferBlendMethod)
};
mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 1);
assert(mBindGroup);
}
void WgBindGroupBlendMethod::release()
{
releaseBuffer(uBufferBlendMethod);
releaseBindGroup(mBindGroup);
}
WGPUBindGroupLayout WgBindGroupCompositeMethod::getLayout(WGPUDevice device)
{
if (layout) return layout;
const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] {
makeBindGroupLayoutEntryBuffer(0)
};
layout = createBindGroupLayout(device, bindGroupLayoutEntries, 1);
assert(layout);
return layout;
}
void WgBindGroupCompositeMethod::releaseLayout()
{
releaseBindGroupLayout(layout);
}
void WgBindGroupCompositeMethod::initialize(WGPUDevice device, WGPUQueue queue, CompositeMethod uCompositeMethod)
{
release();
uint32_t compositeMethod = (uint32_t)uCompositeMethod;
uBufferCompositeMethod = createBuffer(device, queue, &compositeMethod, sizeof(compositeMethod));
const WGPUBindGroupEntry bindGroupEntries[] {
makeBindGroupEntryBuffer(0, uBufferCompositeMethod)
};
mBindGroup = createBindGroup(device, getLayout(device), bindGroupEntries, 1);
assert(mBindGroup);
}
void WgBindGroupCompositeMethod::release()
{
releaseBuffer(uBufferCompositeMethod);
releaseBindGroup(mBindGroup);
}
//************************************************************************ //************************************************************************
// bind group pools // bind group pools
//************************************************************************ //************************************************************************
void WgBindGroupOpacityPool::initialize(WgContext& context) void WgBindGroupOpacityPool::initialize(WgContext& context)
{ {
memset(mPool, 0x00, sizeof(mPool)); release(context);
for (uint32_t opacity = 0; opacity < 256; opacity++) {
mPool[opacity] = new WgBindGroupOpacity;
mPool[opacity]->initialize(context.device, context.queue, opacity);
}
} }
@ -438,11 +514,69 @@ void WgBindGroupOpacityPool::release(WgContext& context)
WgBindGroupOpacity* WgBindGroupOpacityPool::allocate(WgContext& context, uint8_t opacity) WgBindGroupOpacity* WgBindGroupOpacityPool::allocate(WgContext& context, uint8_t opacity)
{ {
WgBindGroupOpacity* bindGroupOpacity = mPool[opacity]; return mPool[opacity];
if (!bindGroupOpacity) { }
bindGroupOpacity = new WgBindGroupOpacity;
bindGroupOpacity->initialize(context.device, context.queue, opacity);
mPool[opacity] = bindGroupOpacity; void WgBindGroupBlendMethodPool::initialize(WgContext& context)
} {
return bindGroupOpacity; release(context);
for (uint8_t blendMethod = (uint8_t)BlendMethod::Normal;
blendMethod <= (uint8_t)BlendMethod::SoftLight;
blendMethod++) {
mPool[blendMethod] = new WgBindGroupBlendMethod;
mPool[blendMethod]->initialize(context.device, context.queue, (BlendMethod)blendMethod);
}
}
void WgBindGroupBlendMethodPool::release(WgContext& context)
{
for (uint8_t blendMethod = (uint8_t)BlendMethod::Normal;
blendMethod <= (uint8_t)BlendMethod::SoftLight;
blendMethod++) {
if (mPool[blendMethod]) {
mPool[blendMethod]->release();
delete mPool[blendMethod];
mPool[blendMethod] = nullptr;
}
}
}
WgBindGroupBlendMethod* WgBindGroupBlendMethodPool::allocate(WgContext& context, BlendMethod blendMethod)
{
return mPool[(uint8_t)blendMethod];
}
void WgBindGroupCompositeMethodPool::initialize(WgContext& context)
{
release(context);
for (uint8_t composeMethos = (uint8_t)CompositeMethod::None;
composeMethos <= (uint8_t)CompositeMethod::DifferenceMask;
composeMethos++) {
mPool[composeMethos] = new WgBindGroupCompositeMethod;
mPool[composeMethos]->initialize(context.device, context.queue, (CompositeMethod)composeMethos);
}
}
void WgBindGroupCompositeMethodPool::release(WgContext& context)
{
for (uint8_t blendMethod = (uint8_t)CompositeMethod::None;
blendMethod <= (uint8_t)CompositeMethod::DifferenceMask;
blendMethod++) {
if (mPool[blendMethod]) {
mPool[blendMethod]->release();
delete mPool[blendMethod];
mPool[blendMethod] = nullptr;
}
}
}
WgBindGroupCompositeMethod* WgBindGroupCompositeMethodPool::allocate(WgContext& context, CompositeMethod composeMethod)
{
return mPool[(uint8_t)composeMethod];
} }

75
src/renderer/wg_engine/tvgWgBindGroups.h
View file

@ -106,19 +106,6 @@ struct WgBindGroupPicture : public WgBindGroup
void release(); void release();
}; };
// @group(1 or 2)
struct WgBindGroupOpacity : public WgBindGroup
{
static WGPUBindGroupLayout layout;
static WGPUBindGroupLayout getLayout(WGPUDevice device);
static void releaseLayout();
WGPUBuffer uBufferOpacity{};
void initialize(WGPUDevice device, WGPUQueue queue, uint32_t uOpacity);
void update(WGPUDevice device, WGPUQueue queue, uint32_t uOpacity);
void release();
};
// @group(0 or 1) // @group(0 or 1)
struct WgBindGroupTexture : public WgBindGroup struct WgBindGroupTexture : public WgBindGroup
{ {
@ -132,7 +119,7 @@ struct WgBindGroupTexture : public WgBindGroup
}; };
// @group(0 or 1) // @group(0 or 1)
struct WgBindGroupStorageTexture : public WgBindGroup struct WgBindGroupTextureStorage : public WgBindGroup
{ {
static WGPUBindGroupLayout layout; static WGPUBindGroupLayout layout;
static WGPUBindGroupLayout getLayout(WGPUDevice device); static WGPUBindGroupLayout getLayout(WGPUDevice device);
@ -156,6 +143,44 @@ struct WgBindGroupTextureSampled : public WgBindGroup
void release(); void release();
}; };
// @group(1 or 2)
struct WgBindGroupOpacity : public WgBindGroup
{
static WGPUBindGroupLayout layout;
static WGPUBindGroupLayout getLayout(WGPUDevice device);
static void releaseLayout();
WGPUBuffer uBufferOpacity{};
void initialize(WGPUDevice device, WGPUQueue queue, uint32_t uOpacity);
void release();
};
// @group(2)
struct WgBindGroupBlendMethod : public WgBindGroup
{
static WGPUBindGroupLayout layout;
static WGPUBindGroupLayout getLayout(WGPUDevice device);
static void releaseLayout();
WGPUBuffer uBufferBlendMethod{};
void initialize(WGPUDevice device, WGPUQueue queue,
BlendMethod uBlendMethod);
void release();
};
// @group(2)
struct WgBindGroupCompositeMethod : public WgBindGroup
{
static WGPUBindGroupLayout layout;
static WGPUBindGroupLayout getLayout(WGPUDevice device);
static void releaseLayout();
WGPUBuffer uBufferCompositeMethod{};
void initialize(WGPUDevice device, WGPUQueue queue,
CompositeMethod uCompositeMethod);
void release();
};
//************************************************************************ //************************************************************************
// bind group pools // bind group pools
//************************************************************************ //************************************************************************
@ -163,11 +188,31 @@ struct WgBindGroupTextureSampled : public WgBindGroup
class WgBindGroupOpacityPool class WgBindGroupOpacityPool
{ {
private: private:
WgBindGroupOpacity* mPool[256]; WgBindGroupOpacity* mPool[256]{};
public: public:
void initialize(WgContext& context); void initialize(WgContext& context);
void release(WgContext& context); void release(WgContext& context);
WgBindGroupOpacity* allocate(WgContext& context, uint8_t opacity); WgBindGroupOpacity* allocate(WgContext& context, uint8_t opacity);
}; };
class WgBindGroupBlendMethodPool
{
private:
WgBindGroupBlendMethod* mPool[(uint8_t)BlendMethod::SoftLight + 1]{};
public:
void initialize(WgContext& context);
void release(WgContext& context);
WgBindGroupBlendMethod* allocate(WgContext& context, BlendMethod blendMethod);
};
class WgBindGroupCompositeMethodPool
{
private:
WgBindGroupCompositeMethod* mPool[(uint8_t)CompositeMethod::DifferenceMask + 1]{};
public:
void initialize(WgContext& context);
void release(WgContext& context);
WgBindGroupCompositeMethod* allocate(WgContext& context, CompositeMethod composeMethod);
};
#endif // _TVG_WG_BIND_GROUPS_H_ #endif // _TVG_WG_BIND_GROUPS_H_

55
src/renderer/wg_engine/tvgWgCommon.cpp
View file

@ -88,11 +88,19 @@ void WgContext::initialize()
queue = wgpuDeviceGetQueue(device); queue = wgpuDeviceGetQueue(device);
assert(queue); assert(queue);
// create default nearest and linear samplers
samplerNearest = createSampler(WGPUFilterMode_Nearest, WGPUMipmapFilterMode_Nearest);
assert(samplerNearest);
samplerLinear = createSampler(WGPUFilterMode_Linear, WGPUMipmapFilterMode_Linear);
assert(samplerLinear);
} }
void WgContext::release() void WgContext::release()
{ {
releaseSampler(samplerNearest);
releaseSampler(samplerLinear);
if (device) { if (device) {
wgpuDeviceDestroy(device); wgpuDeviceDestroy(device);
wgpuDeviceRelease(device); wgpuDeviceRelease(device);
@ -115,7 +123,26 @@ void WgContext::executeCommandEncoder(WGPUCommandEncoder commandEncoder)
} }
WGPUTexture WgContext::createTexture2d(WGPUTextureUsage usage, WGPUTextureFormat format, uint32_t width, uint32_t height, char const * label) { WGPUSampler WgContext::createSampler(WGPUFilterMode minFilter, WGPUMipmapFilterMode mipmapFilter)
{
WGPUSamplerDescriptor samplerDesc{};
samplerDesc.nextInChain = nullptr;
samplerDesc.label = "The sampler";
samplerDesc.addressModeU = WGPUAddressMode_ClampToEdge;
samplerDesc.addressModeV = WGPUAddressMode_ClampToEdge;
samplerDesc.addressModeW = WGPUAddressMode_ClampToEdge;
samplerDesc.magFilter = minFilter;
samplerDesc.minFilter = minFilter;
samplerDesc.mipmapFilter = mipmapFilter;
samplerDesc.lodMinClamp = 0.0f;
samplerDesc.lodMaxClamp = 32.0f;
samplerDesc.compare = WGPUCompareFunction_Undefined;
samplerDesc.maxAnisotropy = 1;
return wgpuDeviceCreateSampler(device, &samplerDesc);
}
WGPUTexture WgContext::createTexture2d(WGPUTextureUsageFlags usage, WGPUTextureFormat format, uint32_t width, uint32_t height, char const * label) {
WGPUTextureDescriptor textureDesc{}; WGPUTextureDescriptor textureDesc{};
textureDesc.nextInChain = nullptr; textureDesc.nextInChain = nullptr;
textureDesc.label = label; textureDesc.label = label;
@ -147,7 +174,16 @@ WGPUTextureView WgContext::createTextureView2d(WGPUTexture texture, WGPU_NULLABL
}; };
void WgContext::releaseTexture(WGPUTexture& texture) { void WgContext::releaseSampler(WGPUSampler& sampler)
{
if (sampler) {
wgpuSamplerRelease(sampler);
sampler = nullptr;
}
}
void WgContext::releaseTexture(WGPUTexture& texture)
{
if (texture) { if (texture) {
wgpuTextureDestroy(texture); wgpuTextureDestroy(texture);
wgpuTextureRelease(texture); wgpuTextureRelease(texture);
@ -157,7 +193,8 @@ void WgContext::releaseTexture(WGPUTexture& texture) {
} }
void WgContext::releaseTextureView(WGPUTextureView& textureView) { void WgContext::releaseTextureView(WGPUTextureView& textureView)
{
if (textureView) wgpuTextureViewRelease(textureView); if (textureView) wgpuTextureViewRelease(textureView);
textureView = nullptr; textureView = nullptr;
} }
@ -226,7 +263,7 @@ WGPUBindGroupLayoutEntry WgBindGroup::makeBindGroupLayoutEntryBuffer(uint32_t bi
WGPUBindGroupLayoutEntry bindGroupLayoutEntry{}; WGPUBindGroupLayoutEntry bindGroupLayoutEntry{};
bindGroupLayoutEntry.nextInChain = nullptr; bindGroupLayoutEntry.nextInChain = nullptr;
bindGroupLayoutEntry.binding = binding; bindGroupLayoutEntry.binding = binding;
bindGroupLayoutEntry.visibility = WGPUShaderStage_Vertex | WGPUShaderStage_Fragment; bindGroupLayoutEntry.visibility = WGPUShaderStage_Vertex | WGPUShaderStage_Fragment | WGPUShaderStage_Compute;
bindGroupLayoutEntry.buffer.nextInChain = nullptr; bindGroupLayoutEntry.buffer.nextInChain = nullptr;
bindGroupLayoutEntry.buffer.type = WGPUBufferBindingType_Uniform; bindGroupLayoutEntry.buffer.type = WGPUBufferBindingType_Uniform;
bindGroupLayoutEntry.buffer.hasDynamicOffset = false; bindGroupLayoutEntry.buffer.hasDynamicOffset = false;
@ -261,14 +298,14 @@ WGPUBindGroupLayoutEntry WgBindGroup::makeBindGroupLayoutEntryTexture(uint32_t b
} }
WGPUBindGroupLayoutEntry WgBindGroup::makeBindGroupLayoutEntryStorageTexture(uint32_t binding) WGPUBindGroupLayoutEntry WgBindGroup::makeBindGroupLayoutEntryStorageTexture(uint32_t binding, WGPUStorageTextureAccess access)
{ {
WGPUBindGroupLayoutEntry bindGroupLayoutEntry{}; WGPUBindGroupLayoutEntry bindGroupLayoutEntry{};
bindGroupLayoutEntry.nextInChain = nullptr; bindGroupLayoutEntry.nextInChain = nullptr;
bindGroupLayoutEntry.binding = binding; bindGroupLayoutEntry.binding = binding;
bindGroupLayoutEntry.visibility = WGPUShaderStage_Fragment | WGPUShaderStage_Compute; bindGroupLayoutEntry.visibility = WGPUShaderStage_Fragment | WGPUShaderStage_Compute;
bindGroupLayoutEntry.storageTexture.nextInChain = nullptr; bindGroupLayoutEntry.storageTexture.nextInChain = nullptr;
bindGroupLayoutEntry.storageTexture.access = WGPUStorageTextureAccess_ReadWrite; bindGroupLayoutEntry.storageTexture.access = access;
bindGroupLayoutEntry.storageTexture.format = WGPUTextureFormat_RGBA8Unorm; bindGroupLayoutEntry.storageTexture.format = WGPUTextureFormat_RGBA8Unorm;
bindGroupLayoutEntry.storageTexture.viewDimension = WGPUTextureViewDimension_2D; bindGroupLayoutEntry.storageTexture.viewDimension = WGPUTextureViewDimension_2D;
return bindGroupLayoutEntry; return bindGroupLayoutEntry;
@ -427,11 +464,11 @@ WGPUBlendState WgRenderPipeline::makeBlendState()
{ {
WGPUBlendState blendState{}; WGPUBlendState blendState{};
blendState.color.operation = WGPUBlendOperation_Add; blendState.color.operation = WGPUBlendOperation_Add;
blendState.color.srcFactor = WGPUBlendFactor_SrcAlpha; blendState.color.srcFactor = WGPUBlendFactor_One;
blendState.color.dstFactor = WGPUBlendFactor_OneMinusSrcAlpha; blendState.color.dstFactor = WGPUBlendFactor_Zero;
blendState.alpha.operation = WGPUBlendOperation_Add; blendState.alpha.operation = WGPUBlendOperation_Add;
blendState.alpha.srcFactor = WGPUBlendFactor_One; blendState.alpha.srcFactor = WGPUBlendFactor_One;
blendState.alpha.dstFactor = WGPUBlendFactor_One; blendState.alpha.dstFactor = WGPUBlendFactor_Zero;
return blendState; return blendState;
} }

12
src/renderer/wg_engine/tvgWgCommon.h
View file

@ -40,6 +40,9 @@ struct WgContext {
WGPUAdapterProperties adapterProperties{}; WGPUAdapterProperties adapterProperties{};
WGPUSupportedLimits supportedLimits{}; WGPUSupportedLimits supportedLimits{};
WGPUSampler samplerNearest{};
WGPUSampler samplerLinear{};
WgPipelines* pipelines{}; // external handle (do not release) WgPipelines* pipelines{}; // external handle (do not release)
void initialize(); void initialize();
@ -47,8 +50,10 @@ struct WgContext {
void executeCommandEncoder(WGPUCommandEncoder commandEncoder); void executeCommandEncoder(WGPUCommandEncoder commandEncoder);
WGPUTexture createTexture2d(WGPUTextureUsage usage, WGPUTextureFormat format, uint32_t width, uint32_t height, char const * label); WGPUSampler createSampler(WGPUFilterMode minFilter, WGPUMipmapFilterMode mipmapFilter);
WGPUTexture createTexture2d(WGPUTextureUsageFlags usage, WGPUTextureFormat format, uint32_t width, uint32_t height, char const * label);
WGPUTextureView createTextureView2d(WGPUTexture texture, WGPU_NULLABLE char const * label); WGPUTextureView createTextureView2d(WGPUTexture texture, WGPU_NULLABLE char const * label);
void releaseSampler(WGPUSampler& sampler);
void releaseTexture(WGPUTexture& texture); void releaseTexture(WGPUTexture& texture);
void releaseTextureView(WGPUTextureView& textureView); void releaseTextureView(WGPUTextureView& textureView);
}; };
@ -67,7 +72,7 @@ struct WgBindGroup
static WGPUBindGroupLayoutEntry makeBindGroupLayoutEntryBuffer(uint32_t binding); static WGPUBindGroupLayoutEntry makeBindGroupLayoutEntryBuffer(uint32_t binding);
static WGPUBindGroupLayoutEntry makeBindGroupLayoutEntrySampler(uint32_t binding); static WGPUBindGroupLayoutEntry makeBindGroupLayoutEntrySampler(uint32_t binding);
static WGPUBindGroupLayoutEntry makeBindGroupLayoutEntryTexture(uint32_t binding); static WGPUBindGroupLayoutEntry makeBindGroupLayoutEntryTexture(uint32_t binding);
static WGPUBindGroupLayoutEntry makeBindGroupLayoutEntryStorageTexture(uint32_t binding); static WGPUBindGroupLayoutEntry makeBindGroupLayoutEntryStorageTexture(uint32_t binding, WGPUStorageTextureAccess access);
static WGPUBuffer createBuffer(WGPUDevice device, WGPUQueue queue, const void *data, size_t size); static WGPUBuffer createBuffer(WGPUDevice device, WGPUQueue queue, const void *data, size_t size);
static WGPUBindGroup createBindGroup(WGPUDevice device, WGPUBindGroupLayout layout, const WGPUBindGroupEntry* bindGroupEntries, uint32_t count); static WGPUBindGroup createBindGroup(WGPUDevice device, WGPUBindGroupLayout layout, const WGPUBindGroupEntry* bindGroupEntries, uint32_t count);
@ -123,6 +128,9 @@ public:
static void destroyRenderPipeline(WGPURenderPipeline& renderPipeline); static void destroyRenderPipeline(WGPURenderPipeline& renderPipeline);
}; };
#define WG_COMPUTE_WORKGROUP_SIZE_X 8
#define WG_COMPUTE_WORKGROUP_SIZE_Y 8
struct WgComputePipeline: public WgPipeline struct WgComputePipeline: public WgPipeline
{ {
protected: protected:

210
src/renderer/wg_engine/tvgWgPipelines.cpp
View file

@ -25,6 +25,10 @@
#define ARRAY_ELEMENTS_COUNT(arr) sizeof(arr)/sizeof(arr[0]) #define ARRAY_ELEMENTS_COUNT(arr) sizeof(arr)/sizeof(arr[0])
//************************************************************************
// graphics pipelines
//************************************************************************
void WgPipelineFillShape::initialize(WGPUDevice device) void WgPipelineFillShape::initialize(WGPUDevice device)
{ {
// vertex attributes settings // vertex attributes settings
@ -222,120 +226,64 @@ void WgPipelineImage::initialize(WGPUDevice device)
shaderSource, shaderLabel, pipelineLabel); shaderSource, shaderLabel, pipelineLabel);
} }
//************************************************************************
// compute pipelines
//************************************************************************
void WgPipelineBlit::initialize(WGPUDevice device) void WgPipelineClear::initialize(WGPUDevice device)
{ {
// vertex and buffers settings
WGPUVertexAttribute vertexAttributesPos = { WGPUVertexFormat_Float32x2, sizeof(float) * 0, 0 };
WGPUVertexAttribute vertexAttributesTex = { WGPUVertexFormat_Float32x2, sizeof(float) * 0, 1 };
WGPUVertexBufferLayout vertexBufferLayouts[] = {
makeVertexBufferLayout(&vertexAttributesPos, 1, sizeof(float) * 2),
makeVertexBufferLayout(&vertexAttributesTex, 1, sizeof(float) * 2)
};
// bind groups and layouts // bind groups and layouts
WGPUBindGroupLayout bindGroupLayouts[] = { WGPUBindGroupLayout bindGroupLayouts[] = {
WgBindGroupTextureSampled::getLayout(device), WgBindGroupTextureStorage::getLayout(device)
};
// sheder source and labels
auto shaderSource = cShaderSource_PipelineComputeClear;
auto shaderLabel = "The compute shader clear";
auto pipelineLabel = "The compute pipeline clear";
// allocate all pipeline handles
allocate(device,
bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts),
shaderSource, shaderLabel, pipelineLabel);
}
void WgPipelineBlend::initialize(WGPUDevice device)
{
// bind groups and layouts
WGPUBindGroupLayout bindGroupLayouts[] = {
WgBindGroupTextureStorage::getLayout(device),
WgBindGroupTextureStorage::getLayout(device),
WgBindGroupBlendMethod::getLayout(device)
};
// sheder source and labels
auto shaderSource = cShaderSource_PipelineComputeBlend;
auto shaderLabel = "The compute shader blend";
auto pipelineLabel = "The compute pipeline blend";
// allocate all pipeline handles
allocate(device,
bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts),
shaderSource, shaderLabel, pipelineLabel);
}
void WgPipelineCompose::initialize(WGPUDevice device)
{
// bind groups and layouts
WGPUBindGroupLayout bindGroupLayouts[] = {
WgBindGroupTextureStorage::getLayout(device),
WgBindGroupTextureStorage::getLayout(device),
WgBindGroupCompositeMethod::getLayout(device),
WgBindGroupOpacity::getLayout(device) WgBindGroupOpacity::getLayout(device)
}; };
// stencil function
WGPUCompareFunction stencilFuncion = WGPUCompareFunction_Always;
WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero;
// sheder source and labels // sheder source and labels
auto shaderSource = cShaderSource_PipelineBlit; auto shaderSource = cShaderSource_PipelineComputeCompose;
auto shaderLabel = "The shader blit"; auto shaderLabel = "The compute shader compose";
auto pipelineLabel = "The render pipeline blit"; auto pipelineLabel = "The compute pipeline compose";
// allocate all pipeline handles
allocate(device,
vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts),
bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts),
stencilFuncion, stencilOperation,
shaderSource, shaderLabel, pipelineLabel);
}
void WgPipelineBlitColor::initialize(WGPUDevice device)
{
// vertex and buffers settings
WGPUVertexAttribute vertexAttributesPos = { WGPUVertexFormat_Float32x2, sizeof(float) * 0, 0 };
WGPUVertexAttribute vertexAttributesTex = { WGPUVertexFormat_Float32x2, sizeof(float) * 0, 1 };
WGPUVertexBufferLayout vertexBufferLayouts[] = {
makeVertexBufferLayout(&vertexAttributesPos, 1, sizeof(float) * 2),
makeVertexBufferLayout(&vertexAttributesTex, 1, sizeof(float) * 2)
};
// bind groups and layouts
WGPUBindGroupLayout bindGroupLayouts[] = {
WgBindGroupTextureSampled::getLayout(device)
};
// stencil function
WGPUCompareFunction stencilFuncion = WGPUCompareFunction_Always;
WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero;
// sheder source and labels
auto shaderSource = cShaderSource_PipelineBlitColor;
auto shaderLabel = "The shader blit color";
auto pipelineLabel = "The render pipeline blit color";
// allocate all pipeline handles
allocate(device,
vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts),
bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts),
stencilFuncion, stencilOperation,
shaderSource, shaderLabel, pipelineLabel);
}
void WgPipelineComposition::initialize(WGPUDevice device, const char* shaderSrc)
{
// vertex and buffers settings
WGPUVertexAttribute vertexAttributesPos = { WGPUVertexFormat_Float32x2, sizeof(float) * 0, 0 };
WGPUVertexAttribute vertexAttributesTex = { WGPUVertexFormat_Float32x2, sizeof(float) * 0, 1 };
WGPUVertexBufferLayout vertexBufferLayouts[] = {
makeVertexBufferLayout(&vertexAttributesPos, 1, sizeof(float) * 2),
makeVertexBufferLayout(&vertexAttributesTex, 1, sizeof(float) * 2)
};
// bind groups and layouts
WGPUBindGroupLayout bindGroupLayouts[] = {
WgBindGroupTextureSampled::getLayout(device),
WgBindGroupTextureSampled::getLayout(device)
};
// stencil function
WGPUCompareFunction stencilFuncion = WGPUCompareFunction_Always;
WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero;
// sheder source and labels
auto shaderSource = shaderSrc;
auto shaderLabel = "The shader compose alpha mask";
auto pipelineLabel = "The render pipeline compose alpha mask";
// allocate all pipeline handles
allocate(device,
vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts),
bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts),
stencilFuncion, stencilOperation,
shaderSource, shaderLabel, pipelineLabel);
}
void WgPipelineBlend::initialize(WGPUDevice device, const char* shaderSrc)
{
// bind groups and layouts
WGPUBindGroupLayout bindGroupLayouts[] = {
//WgBindGroupTexture::getLayout(device),
WgBindGroupStorageTexture::getLayout(device),
WgBindGroupStorageTexture::getLayout(device)
};
// sheder source and labels
auto shaderSource = shaderSrc;
auto shaderLabel = "The compute shader blend";
auto pipelineLabel = "The compute pipeline blend";
// allocate all pipeline handles // allocate all pipeline handles
allocate(device, allocate(device,
@ -356,20 +304,10 @@ void WgPipelines::initialize(WgContext& context)
linear.initialize(context.device); linear.initialize(context.device);
radial.initialize(context.device); radial.initialize(context.device);
image.initialize(context.device); image.initialize(context.device);
// blit pipelines
blit.initialize(context.device);
blitColor.initialize(context.device);
// composition pipelines
compAlphaMask.initialize(context.device, cShaderSource_PipelineCompAlphaMask);
compInvAlphaMask.initialize(context.device, cShaderSource_PipelineCompInvAlphaMask);
compLumaMask.initialize(context.device, cShaderSource_PipelineCompLumaMask);
compInvLumaMask.initialize(context.device, cShaderSource_PipelineCompInvLumaMask);
compAddMask.initialize(context.device, cShaderSource_PipelineCompAddMask);
compSubtractMask.initialize(context.device, cShaderSource_PipelineCompSubtractMask);
compIntersectMask.initialize(context.device, cShaderSource_PipelineCompIntersectMask);
compDifferenceMask.initialize(context.device, cShaderSource_PipelineCompDifferenceMask);
// compute pipelines // compute pipelines
computeBlend.initialize(context.device, cShaderSource_PipelineComputeBlend); computeClear.initialize(context.device);
computeBlend.initialize(context.device);
computeCompose.initialize(context.device);
// store pipelines to context // store pipelines to context
context.pipelines = this; context.pipelines = this;
} }
@ -378,7 +316,7 @@ void WgPipelines::initialize(WgContext& context)
void WgPipelines::release() void WgPipelines::release()
{ {
WgBindGroupTextureSampled::releaseLayout(); WgBindGroupTextureSampled::releaseLayout();
WgBindGroupStorageTexture::releaseLayout(); WgBindGroupTextureStorage::releaseLayout();
WgBindGroupTexture::releaseLayout(); WgBindGroupTexture::releaseLayout();
WgBindGroupOpacity::releaseLayout(); WgBindGroupOpacity::releaseLayout();
WgBindGroupPicture::releaseLayout(); WgBindGroupPicture::releaseLayout();
@ -388,19 +326,9 @@ void WgPipelines::release()
WgBindGroupPaint::releaseLayout(); WgBindGroupPaint::releaseLayout();
WgBindGroupCanvas::releaseLayout(); WgBindGroupCanvas::releaseLayout();
// compute pipelines // compute pipelines
computeCompose.release();
computeBlend.release(); computeBlend.release();
// composition pipelines computeClear.release();
compDifferenceMask.release();
compIntersectMask.release();
compSubtractMask.release();
compAddMask.release();
compInvLumaMask.release();
compLumaMask.release();
compInvAlphaMask.release();
compAlphaMask.release();
// blit pipelines
blitColor.release();
blit.release();
// fill pipelines // fill pipelines
image.release(); image.release();
radial.release(); radial.release();
@ -409,21 +337,3 @@ void WgPipelines::release()
fillStroke.release(); fillStroke.release();
fillShape.release(); fillShape.release();
} }
WgPipelineComposition* WgPipelines::getCompositionPipeline(CompositeMethod method)
{
switch (method) {
case CompositeMethod::ClipPath:
case CompositeMethod::AlphaMask: return &compAlphaMask; break;
case CompositeMethod::InvAlphaMask: return &compInvAlphaMask; break;
case CompositeMethod::LumaMask: return &compLumaMask; break;
case CompositeMethod::InvLumaMask: return &compInvLumaMask; break;
case CompositeMethod::AddMask: return &compAddMask; break;
case CompositeMethod::SubtractMask: return &compSubtractMask; break;
case CompositeMethod::IntersectMask: return &compIntersectMask; break;
case CompositeMethod::DifferenceMask: return &compDifferenceMask; break;
default: return nullptr; break;
}
return nullptr;
}

93
src/renderer/wg_engine/tvgWgPipelines.h
View file

@ -99,59 +99,52 @@ struct WgPipelineImage: public WgRenderPipeline
} }
}; };
struct WgPipelineBlit: public WgRenderPipeline
{
void initialize(WGPUDevice device) override;
void use(WGPURenderPassEncoder encoder,
WgBindGroupTextureSampled& groupTexSampled,
WgBindGroupOpacity& groupOpacity)
{
set(encoder);
groupTexSampled.set(encoder, 0);
groupOpacity.set(encoder, 1);
}
};
struct WgPipelineBlitColor: public WgRenderPipeline
{
void initialize(WGPUDevice device) override;
void use(WGPURenderPassEncoder encoder,
WgBindGroupTextureSampled& groupTexSampled)
{
set(encoder);
groupTexSampled.set(encoder, 0);
}
};
struct WgPipelineComposition: public WgRenderPipeline
{
void initialize(WGPUDevice device) override {};
void initialize(WGPUDevice device, const char* shaderSrc);
void use(WGPURenderPassEncoder encoder,
WgBindGroupTextureSampled& groupTexSampledSrc,
WgBindGroupTextureSampled& groupTexSampledMsk)
{
set(encoder);
groupTexSampledSrc.set(encoder, 0);
groupTexSampledMsk.set(encoder, 1);
}
};
//***************************************************************************** //*****************************************************************************
// compute pipelines // compute pipelines
//***************************************************************************** //*****************************************************************************
struct WgPipelineClear: public WgComputePipeline
{
void initialize(WGPUDevice device) override;
void use(WGPUComputePassEncoder encoder,
WgBindGroupTextureStorage& groupTexDst)
{
set(encoder);
groupTexDst.set(encoder, 0);
}
};
struct WgPipelineBlend: public WgComputePipeline struct WgPipelineBlend: public WgComputePipeline
{ {
void initialize(WGPUDevice device) override {}; void initialize(WGPUDevice device) override;
void initialize(WGPUDevice device, const char* shaderSrc);
void use(WGPUComputePassEncoder encoder, void use(WGPUComputePassEncoder encoder,
WgBindGroupStorageTexture& groupTexSrc, WgBindGroupTextureStorage& groupTexSrc,
WgBindGroupStorageTexture& groupTexDst) WgBindGroupTextureStorage& groupTexDst,
WgBindGroupBlendMethod& blendMethod)
{ {
set(encoder); set(encoder);
groupTexSrc.set(encoder, 0); groupTexSrc.set(encoder, 0);
groupTexDst.set(encoder, 1); groupTexDst.set(encoder, 1);
blendMethod.set(encoder, 2);
}
};
struct WgPipelineCompose: public WgComputePipeline
{
void initialize(WGPUDevice device) override;
void use(WGPUComputePassEncoder encoder,
WgBindGroupTextureStorage& groupTexSrc,
WgBindGroupTextureStorage& groupTexMsk,
WgBindGroupCompositeMethod& groupComposeMethod,
WgBindGroupOpacity& groupOpacity)
{
set(encoder);
groupTexSrc.set(encoder, 0);
groupTexMsk.set(encoder, 1);
groupComposeMethod.set(encoder, 2);
groupOpacity.set(encoder, 3);
} }
}; };
@ -161,30 +154,20 @@ struct WgPipelineBlend: public WgComputePipeline
struct WgPipelines struct WgPipelines
{ {
// render pipelines
WgPipelineFillShape fillShape; WgPipelineFillShape fillShape;
WgPipelineFillStroke fillStroke; WgPipelineFillStroke fillStroke;
WgPipelineSolid solid; WgPipelineSolid solid;
WgPipelineLinear linear; WgPipelineLinear linear;
WgPipelineRadial radial; WgPipelineRadial radial;
WgPipelineImage image; WgPipelineImage image;
WgPipelineBlit blit;
WgPipelineBlitColor blitColor;
// composition pipelines
WgPipelineComposition compAlphaMask;
WgPipelineComposition compInvAlphaMask;
WgPipelineComposition compLumaMask;
WgPipelineComposition compInvLumaMask;
WgPipelineComposition compAddMask;
WgPipelineComposition compSubtractMask;
WgPipelineComposition compIntersectMask;
WgPipelineComposition compDifferenceMask;
// compute pipelines // compute pipelines
WgPipelineClear computeClear;
WgPipelineBlend computeBlend; WgPipelineBlend computeBlend;
WgPipelineCompose computeCompose;
void initialize(WgContext& context); void initialize(WgContext& context);
void release(); void release();
WgPipelineComposition* getCompositionPipeline(CompositeMethod method);
}; };
#endif // _TVG_WG_PIPELINES_H_ #endif // _TVG_WG_PIPELINES_H_

61
src/renderer/wg_engine/tvgWgRenderData.cpp
View file

@ -144,48 +144,12 @@ void WgImageData::update(WgContext& context, Surface* surface)
{ {
release(context); release(context);
assert(surface); assert(surface);
// sampler descriptor texture = context.createTexture2d(
WGPUSamplerDescriptor samplerDesc{}; WGPUTextureUsage_TextureBinding | WGPUTextureUsage_CopyDst,
samplerDesc.nextInChain = nullptr; WGPUTextureFormat_RGBA8Unorm,
samplerDesc.label = "The shape sampler"; surface->w, surface->h, "The shape texture");
samplerDesc.addressModeU = WGPUAddressMode_ClampToEdge;
samplerDesc.addressModeV = WGPUAddressMode_ClampToEdge;
samplerDesc.addressModeW = WGPUAddressMode_ClampToEdge;
samplerDesc.magFilter = WGPUFilterMode_Nearest;
samplerDesc.minFilter = WGPUFilterMode_Nearest;
samplerDesc.mipmapFilter = WGPUMipmapFilterMode_Nearest;
samplerDesc.lodMinClamp = 0.0f;
samplerDesc.lodMaxClamp = 32.0f;
samplerDesc.compare = WGPUCompareFunction_Undefined;
samplerDesc.maxAnisotropy = 1;
sampler = wgpuDeviceCreateSampler(context.device, &samplerDesc);
assert(sampler);
// texture descriptor
WGPUTextureDescriptor textureDesc{};
textureDesc.nextInChain = nullptr;
textureDesc.label = "The shape texture";
textureDesc.usage = WGPUTextureUsage_TextureBinding | WGPUTextureUsage_CopyDst;
textureDesc.dimension = WGPUTextureDimension_2D;
textureDesc.size = { surface->w, surface->h, 1 };
textureDesc.format = WGPUTextureFormat_RGBA8Unorm;
textureDesc.mipLevelCount = 1;
textureDesc.sampleCount = 1;
textureDesc.viewFormatCount = 0;
textureDesc.viewFormats = nullptr;
texture = wgpuDeviceCreateTexture(context.device, &textureDesc);
assert(texture); assert(texture);
// texture view descriptor textureView = context.createTextureView2d(texture, "The shape texture view");
WGPUTextureViewDescriptor textureViewDesc{};
textureViewDesc.nextInChain = nullptr;
textureViewDesc.label = "The shape texture view";
textureViewDesc.format = WGPUTextureFormat_RGBA8Unorm;
textureViewDesc.dimension = WGPUTextureViewDimension_2D;
textureViewDesc.baseMipLevel = 0;
textureViewDesc.mipLevelCount = 1;
textureViewDesc.baseArrayLayer = 0;
textureViewDesc.arrayLayerCount = 1;
textureViewDesc.aspect = WGPUTextureAspect_All;
textureView = wgpuTextureCreateView(texture, &textureViewDesc);
assert(textureView); assert(textureView);
// update texture data // update texture data
WGPUImageCopyTexture imageCopyTexture{}; WGPUImageCopyTexture imageCopyTexture{};
@ -209,19 +173,8 @@ void WgImageData::update(WgContext& context, Surface* surface)
void WgImageData::release(WgContext& context) void WgImageData::release(WgContext& context)
{ {
if (textureView) { context.releaseTextureView(textureView);
wgpuTextureViewRelease(textureView); context.releaseTexture(texture);
textureView = nullptr;
}
if (texture) {
wgpuTextureDestroy(texture);
wgpuTextureRelease(texture);
texture = nullptr;
}
if (sampler) {
wgpuSamplerRelease(sampler);
sampler = nullptr;
}
}; };
//*********************************************************************** //***********************************************************************

1
src/renderer/wg_engine/tvgWgRenderData.h
View file

@ -45,7 +45,6 @@ struct WgMeshDataGroup {
}; };
struct WgImageData { struct WgImageData {
WGPUSampler sampler{};
WGPUTexture texture{}; WGPUTexture texture{};
WGPUTextureView textureView{}; WGPUTextureView textureView{};

436
src/renderer/wg_engine/tvgWgRenderTarget.cpp
View file

@ -22,133 +22,128 @@
#include "tvgWgRenderTarget.h" #include "tvgWgRenderTarget.h"
//*****************************************************************************
// render target
//*****************************************************************************
void WgRenderTarget::initialize(WgContext& context, uint32_t w, uint32_t h) void WgRenderTarget::initialize(WgContext& context, uint32_t w, uint32_t h)
{ {
release(context); release(context);
// sampler descriptor // create color and stencil textures
WGPUSamplerDescriptor samplerDesc{}; texColor = context.createTexture2d(
samplerDesc.nextInChain = nullptr; WGPUTextureUsage_RenderAttachment | WGPUTextureUsage_StorageBinding,
samplerDesc.label = "The target sampler"; WGPUTextureFormat_RGBA8Unorm,
samplerDesc.addressModeU = WGPUAddressMode_ClampToEdge; w, h, "The target texture color");
samplerDesc.addressModeV = WGPUAddressMode_ClampToEdge; texStencil = context.createTexture2d(
samplerDesc.addressModeW = WGPUAddressMode_ClampToEdge; WGPUTextureUsage_RenderAttachment,
samplerDesc.magFilter = WGPUFilterMode_Nearest; WGPUTextureFormat_Stencil8,
samplerDesc.minFilter = WGPUFilterMode_Nearest; w, h, "The target texture stencil");
samplerDesc.mipmapFilter = WGPUMipmapFilterMode_Nearest; assert(texColor);
samplerDesc.lodMinClamp = 0.0f; assert(texStencil);
samplerDesc.lodMaxClamp = 32.0f; texViewColor = context.createTextureView2d(texColor, "The target texture view color");
samplerDesc.compare = WGPUCompareFunction_Undefined; texViewStencil = context.createTextureView2d(texStencil, "The target texture view stencil");
samplerDesc.maxAnisotropy = 1; assert(texViewColor);
sampler = wgpuDeviceCreateSampler(context.device, &samplerDesc); assert(texViewStencil);
assert(sampler);
// texture descriptor
WGPUTextureDescriptor textureDescColor{};
textureDescColor.nextInChain = nullptr;
textureDescColor.label = "The target texture color";
textureDescColor.usage = WGPUTextureUsage_RenderAttachment | WGPUTextureUsage_TextureBinding | WGPUTextureUsage_CopyDst | WGPUTextureUsage_StorageBinding;
textureDescColor.dimension = WGPUTextureDimension_2D;
textureDescColor.size = { w, h, 1 };
textureDescColor.format = WGPUTextureFormat_RGBA8Unorm;
textureDescColor.mipLevelCount = 1;
textureDescColor.sampleCount = 1;
textureDescColor.viewFormatCount = 0;
textureDescColor.viewFormats = nullptr;
mTextureColor = wgpuDeviceCreateTexture(context.device, &textureDescColor);
assert(mTextureColor);
// texture view descriptor
WGPUTextureViewDescriptor textureViewDescColor{};
textureViewDescColor.nextInChain = nullptr;
textureViewDescColor.label = "The target texture view color";
textureViewDescColor.format = WGPUTextureFormat_RGBA8Unorm;
textureViewDescColor.dimension = WGPUTextureViewDimension_2D;
textureViewDescColor.baseMipLevel = 0;
textureViewDescColor.mipLevelCount = 1;
textureViewDescColor.baseArrayLayer = 0;
textureViewDescColor.arrayLayerCount = 1;
textureViewDescColor.aspect = WGPUTextureAspect_All;
textureViewColor = wgpuTextureCreateView(mTextureColor, &textureViewDescColor);
assert(textureViewColor);
// stencil texture
WGPUTextureDescriptor textureDescStencil{};
textureDescStencil.nextInChain = nullptr;
textureDescStencil.label = "The target texture stencil";
textureDescStencil.usage = WGPUTextureUsage_RenderAttachment;
textureDescStencil.dimension = WGPUTextureDimension_2D;
textureDescStencil.size = { w, h, 1 };
textureDescStencil.format = WGPUTextureFormat_Stencil8;
textureDescStencil.mipLevelCount = 1;
textureDescStencil.sampleCount = 1;
textureDescStencil.viewFormatCount = 0;
textureDescStencil.viewFormats = nullptr;
mTextureStencil = wgpuDeviceCreateTexture(context.device, &textureDescStencil);
assert(mTextureStencil);
// texture view descriptor
WGPUTextureViewDescriptor textureViewDescStencil{};
textureViewDescStencil.nextInChain = nullptr;
textureViewDescStencil.label = "The target texture view stncil";
textureViewDescStencil.format = WGPUTextureFormat_Stencil8;
textureViewDescStencil.dimension = WGPUTextureViewDimension_2D;
textureViewDescStencil.baseMipLevel = 0;
textureViewDescStencil.mipLevelCount = 1;
textureViewDescStencil.baseArrayLayer = 0;
textureViewDescStencil.arrayLayerCount = 1;
textureViewDescStencil.aspect = WGPUTextureAspect_All;
textureViewStencil = wgpuTextureCreateView(mTextureStencil, &textureViewDescStencil);
assert(textureViewStencil);
// initialize bind group for blitting // initialize bind group for blitting
bindGroupTex.initialize(context.device, context.queue, textureViewColor); bindGroupTexStorage.initialize(context.device, context.queue, texViewColor);
bindGroupStorageTex.initialize(context.device, context.queue, textureViewColor);
bindGroupTexSampled.initialize(context.device, context.queue, sampler, textureViewColor);
// initialize window binding groups // initialize window binding groups
WgShaderTypeMat4x4f viewMat(w, h); WgShaderTypeMat4x4f viewMat(w, h);
mBindGroupCanvasWnd.initialize(context.device, context.queue, viewMat); mBindGroupCanvas.initialize(context.device, context.queue, viewMat);
WgGeometryData geometryDataWnd;
geometryDataWnd.appendBlitBox();
mMeshDataCanvasWnd.update(context, &geometryDataWnd);
mPipelines = context.pipelines; mPipelines = context.pipelines;
} }
void WgRenderTarget::release(WgContext& context) void WgRenderTarget::release(WgContext& context)
{ {
mMeshDataCanvasWnd.release(context); mBindGroupCanvas.release();
mBindGroupCanvasWnd.release(); bindGroupTexStorage.release();
bindGroupTexSampled.release(); context.releaseTextureView(texViewStencil);
bindGroupStorageTex.release(); context.releaseTextureView(texViewColor);
bindGroupTex.release(); context.releaseTexture(texStencil);
if (mTextureStencil) { context.releaseTexture(texColor);
wgpuTextureDestroy(mTextureStencil);
wgpuTextureRelease(mTextureStencil);
mTextureStencil = nullptr;
}
if (textureViewStencil) wgpuTextureViewRelease(textureViewStencil);
textureViewStencil = nullptr;
if (mTextureColor) {
wgpuTextureDestroy(mTextureColor);
wgpuTextureRelease(mTextureColor);
mTextureColor = nullptr;
}
if (textureViewColor) wgpuTextureViewRelease(textureViewColor);
textureViewColor = nullptr;
if (sampler) wgpuSamplerRelease(sampler);
sampler = nullptr;
} }
void WgRenderTarget::beginRenderPass(WGPUCommandEncoder commandEncoder, bool clear) void WgRenderTarget::renderShape(WGPUCommandEncoder commandEncoder, WgRenderDataShape* renderData)
{ {
beginRenderPass(commandEncoder, textureViewColor, clear); assert(renderData);
assert(commandEncoder);
WGPURenderPassEncoder renderPassEncoder = beginRenderPass(commandEncoder);
drawShape(renderPassEncoder, renderData);
drawStroke(renderPassEncoder, renderData);
endRenderPass(renderPassEncoder);
} }
void WgRenderTarget::beginRenderPass(WGPUCommandEncoder commandEncoder, WGPUTextureView colorAttachement, bool clear) void WgRenderTarget::renderPicture(WGPUCommandEncoder commandEncoder, WgRenderDataPicture* renderData)
{ {
assert(!mRenderPassEncoder); assert(renderData);
assert(commandEncoder);
WGPURenderPassEncoder renderPassEncoder = beginRenderPass(commandEncoder);
wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 0);
mPipelines->image.use(renderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint, renderData->bindGroupPicture);
renderData->meshData.drawImage(renderPassEncoder);
endRenderPass(renderPassEncoder);
}
void WgRenderTarget::drawShape(WGPURenderPassEncoder renderPassEncoder, WgRenderDataShape* renderData)
{
assert(renderData);
assert(renderPassEncoder);
// draw shape geometry
wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 0);
for (uint32_t i = 0; i < renderData->meshGroupShapes.meshes.count; i++) {
// draw to stencil (first pass)
mPipelines->fillShape.use(renderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint);
renderData->meshGroupShapes.meshes[i]->draw(renderPassEncoder);
// fill shape (second pass)
WgRenderSettings& settings = renderData->renderSettingsShape;
if (settings.fillType == WgRenderSettingsType::Solid)
mPipelines->solid.use(renderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint, settings.bindGroupSolid);
else if (settings.fillType == WgRenderSettingsType::Linear)
mPipelines->linear.use(renderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint, settings.bindGroupLinear);
else if (settings.fillType == WgRenderSettingsType::Radial)
mPipelines->radial.use(renderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint, settings.bindGroupRadial);
renderData->meshBBoxShapes.draw(renderPassEncoder);
}
}
void WgRenderTarget::drawStroke(WGPURenderPassEncoder renderPassEncoder, WgRenderDataShape* renderData)
{
assert(renderData);
assert(renderPassEncoder);
// draw stroke geometry
if (renderData->meshGroupStrokes.meshes.count > 0) {
// draw strokes to stencil (first pass)
wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 255);
for (uint32_t i = 0; i < renderData->meshGroupStrokes.meshes.count; i++) {
mPipelines->fillStroke.use(renderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint);
renderData->meshGroupStrokes.meshes[i]->draw(renderPassEncoder);
}
// fill shape (second pass)
wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 0);
WgRenderSettings& settings = renderData->renderSettingsStroke;
if (settings.fillType == WgRenderSettingsType::Solid)
mPipelines->solid.use(renderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint, settings.bindGroupSolid);
else if (settings.fillType == WgRenderSettingsType::Linear)
mPipelines->linear.use(renderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint, settings.bindGroupLinear);
else if (settings.fillType == WgRenderSettingsType::Radial)
mPipelines->radial.use(renderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint, settings.bindGroupRadial);
renderData->meshBBoxStrokes.draw(renderPassEncoder);
}
}
WGPURenderPassEncoder WgRenderTarget::beginRenderPass(WGPUCommandEncoder commandEncoder)
{
assert(commandEncoder);
// render pass depth stencil attachment // render pass depth stencil attachment
WGPURenderPassDepthStencilAttachment depthStencilAttachment{}; WGPURenderPassDepthStencilAttachment depthStencilAttachment{};
depthStencilAttachment.view = textureViewStencil; depthStencilAttachment.view = texViewStencil;
depthStencilAttachment.depthLoadOp = WGPULoadOp_Clear; depthStencilAttachment.depthLoadOp = WGPULoadOp_Load;
depthStencilAttachment.depthStoreOp = WGPUStoreOp_Store; depthStencilAttachment.depthStoreOp = WGPUStoreOp_Discard;
depthStencilAttachment.depthClearValue = 1.0f; depthStencilAttachment.depthClearValue = 1.0f;
depthStencilAttachment.depthReadOnly = false; depthStencilAttachment.depthReadOnly = false;
depthStencilAttachment.stencilLoadOp = WGPULoadOp_Clear; depthStencilAttachment.stencilLoadOp = WGPULoadOp_Clear;
@ -157,9 +152,9 @@ void WgRenderTarget::beginRenderPass(WGPUCommandEncoder commandEncoder, WGPUText
depthStencilAttachment.stencilReadOnly = false; depthStencilAttachment.stencilReadOnly = false;
// render pass color attachment // render pass color attachment
WGPURenderPassColorAttachment colorAttachment{}; WGPURenderPassColorAttachment colorAttachment{};
colorAttachment.view = colorAttachement; colorAttachment.view = texViewColor;
colorAttachment.resolveTarget = nullptr; colorAttachment.resolveTarget = nullptr;
colorAttachment.loadOp = clear ? WGPULoadOp_Clear : WGPULoadOp_Load; colorAttachment.loadOp = WGPULoadOp_Clear;
colorAttachment.clearValue = { 0, 0, 0, 0 }; colorAttachment.clearValue = { 0, 0, 0, 0 };
colorAttachment.storeOp = WGPUStoreOp_Store; colorAttachment.storeOp = WGPUStoreOp_Store;
// render pass descriptor // render pass descriptor
@ -173,134 +168,145 @@ void WgRenderTarget::beginRenderPass(WGPUCommandEncoder commandEncoder, WGPUText
renderPassDesc.occlusionQuerySet = nullptr; renderPassDesc.occlusionQuerySet = nullptr;
renderPassDesc.timestampWrites = nullptr; renderPassDesc.timestampWrites = nullptr;
// begin render pass // begin render pass
mRenderPassEncoder = wgpuCommandEncoderBeginRenderPass(commandEncoder, &renderPassDesc); return wgpuCommandEncoderBeginRenderPass(commandEncoder, &renderPassDesc);
} }
void WgRenderTarget::endRenderPass() void WgRenderTarget::endRenderPass(WGPURenderPassEncoder renderPassEncoder)
{ {
assert(mRenderPassEncoder); assert(renderPassEncoder);
wgpuRenderPassEncoderEnd(mRenderPassEncoder); wgpuRenderPassEncoderEnd(renderPassEncoder);
wgpuRenderPassEncoderRelease(mRenderPassEncoder); wgpuRenderPassEncoderRelease(renderPassEncoder);
mRenderPassEncoder = nullptr;
}
void WgRenderTarget::renderShape(WgRenderDataShape* renderData)
{
assert(renderData);
assert(mRenderPassEncoder);
// draw shape geometry
wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 0);
for (uint32_t i = 0; i < renderData->meshGroupShapes.meshes.count; i++) {
// draw to stencil (first pass)
mPipelines->fillShape.use(mRenderPassEncoder, mBindGroupCanvasWnd, renderData->bindGroupPaint);
renderData->meshGroupShapes.meshes[i]->draw(mRenderPassEncoder);
// fill shape (second pass)
WgRenderSettings& settings = renderData->renderSettingsShape;
if (settings.fillType == WgRenderSettingsType::Solid)
mPipelines->solid.use(mRenderPassEncoder, mBindGroupCanvasWnd, renderData->bindGroupPaint, settings.bindGroupSolid);
else if (settings.fillType == WgRenderSettingsType::Linear)
mPipelines->linear.use(mRenderPassEncoder, mBindGroupCanvasWnd, renderData->bindGroupPaint, settings.bindGroupLinear);
else if (settings.fillType == WgRenderSettingsType::Radial)
mPipelines->radial.use(mRenderPassEncoder, mBindGroupCanvasWnd, renderData->bindGroupPaint, settings.bindGroupRadial);
renderData->meshBBoxShapes.draw(mRenderPassEncoder);
}
}
void WgRenderTarget::renderStroke(WgRenderDataShape* renderData)
{
assert(renderData);
assert(mRenderPassEncoder);
// draw stroke geometry
if (renderData->meshGroupStrokes.meshes.count > 0) {
// draw strokes to stencil (first pass)
wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 255);
for (uint32_t i = 0; i < renderData->meshGroupStrokes.meshes.count; i++) {
mPipelines->fillStroke.use(mRenderPassEncoder, mBindGroupCanvasWnd, renderData->bindGroupPaint);
renderData->meshGroupStrokes.meshes[i]->draw(mRenderPassEncoder);
}
// fill shape (second pass)
wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 0);
WgRenderSettings& settings = renderData->renderSettingsStroke;
if (settings.fillType == WgRenderSettingsType::Solid)
mPipelines->solid.use(mRenderPassEncoder, mBindGroupCanvasWnd, renderData->bindGroupPaint, settings.bindGroupSolid);
else if (settings.fillType == WgRenderSettingsType::Linear)
mPipelines->linear.use(mRenderPassEncoder, mBindGroupCanvasWnd, renderData->bindGroupPaint, settings.bindGroupLinear);
else if (settings.fillType == WgRenderSettingsType::Radial)
mPipelines->radial.use(mRenderPassEncoder, mBindGroupCanvasWnd, renderData->bindGroupPaint, settings.bindGroupRadial);
renderData->meshBBoxStrokes.draw(mRenderPassEncoder);
}
}
void WgRenderTarget::renderPicture(WgRenderDataPicture* renderData)
{
assert(renderData);
assert(mRenderPassEncoder);
if (renderData->meshData.bufferTexCoord) {
wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 0);
mPipelines->image.use(
mRenderPassEncoder,
mBindGroupCanvasWnd,
renderData->bindGroupPaint,
renderData->bindGroupPicture);
renderData->meshData.drawImage(mRenderPassEncoder);
}
}
void WgRenderTarget::blit(WgContext& context, WgRenderTarget* renderTargetSrc, WgBindGroupOpacity* mBindGroupOpacity)
{
assert(mRenderPassEncoder);
mPipelines->blit.use(mRenderPassEncoder, renderTargetSrc->bindGroupTexSampled, *mBindGroupOpacity);
mMeshDataCanvasWnd.drawImage(mRenderPassEncoder);
}
void WgRenderTarget::blitColor(WgContext& context, WgRenderTarget* renderTargetSrc)
{
assert(mRenderPassEncoder);
mPipelines->blitColor.use(mRenderPassEncoder, renderTargetSrc->bindGroupTexSampled);
mMeshDataCanvasWnd.drawImage(mRenderPassEncoder);
}
void WgRenderTarget::compose(WgContext& context, WgRenderTarget* renderTargetSrc, WgRenderTarget* renderTargetMsk, CompositeMethod method)
{
assert(mRenderPassEncoder);
WgPipelineComposition* pipeline = mPipelines->getCompositionPipeline(method);
assert(pipeline);
pipeline->use(mRenderPassEncoder, renderTargetSrc->bindGroupTexSampled, renderTargetMsk->bindGroupTexSampled);
mMeshDataCanvasWnd.drawImage(mRenderPassEncoder);
} }
//***************************************************************************** //*****************************************************************************
// render terget pool // render storage
//***************************************************************************** //*****************************************************************************
WgRenderTarget* WgRenderTargetPool::allocate(WgContext& context, uint32_t w, uint32_t h) void WgRenderStorage::initialize(WgContext& context, uint32_t w, uint32_t h)
{
release(context);
// store target storage size
width = w;
height = h;
workgroupsCountX = (width + WG_COMPUTE_WORKGROUP_SIZE_X - 1) / WG_COMPUTE_WORKGROUP_SIZE_X; // workgroup size x == 8
workgroupsCountY = (height + WG_COMPUTE_WORKGROUP_SIZE_Y - 1) / WG_COMPUTE_WORKGROUP_SIZE_Y; // workgroup size y == 8
// create color and stencil textures
texStorage = context.createTexture2d(
WGPUTextureUsage_StorageBinding | WGPUTextureUsage_CopySrc,
WGPUTextureFormat_RGBA8Unorm,
w, h, "The target texture storage color");
assert(texStorage);
texViewStorage = context.createTextureView2d(texStorage, "The target texture storage view color");
assert(texViewStorage);
// initialize bind group for blitting
bindGroupTexStorage.initialize(context.device, context.queue, texViewStorage);
mPipelines = context.pipelines;
}
void WgRenderStorage::release(WgContext& context)
{ {
WgRenderTarget* renderTarget{}; bindGroupTexStorage.release();
context.releaseTextureView(texViewStorage);
context.releaseTexture(texStorage);
workgroupsCountX = 0;
workgroupsCountY = 0;
height = 0;
width = 0;
}
void WgRenderStorage::clear(WGPUCommandEncoder commandEncoder)
{
assert(commandEncoder);
WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder);
mPipelines->computeClear.use(computePassEncoder, bindGroupTexStorage);
dispatchWorkgroups(computePassEncoder);
endRenderPass(computePassEncoder);
}
void WgRenderStorage::blend(WGPUCommandEncoder commandEncoder, WgRenderTarget* targetSrc, WgBindGroupBlendMethod* blendMethod)
{
assert(commandEncoder);
assert(targetSrc);
WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder);
mPipelines->computeBlend.use(computePassEncoder, targetSrc->bindGroupTexStorage, bindGroupTexStorage, *blendMethod);
dispatchWorkgroups(computePassEncoder);
endRenderPass(computePassEncoder);
};
void WgRenderStorage::blend(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc, WgBindGroupBlendMethod* blendMethod)
{
assert(commandEncoder);
assert(targetSrc);
WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder);
mPipelines->computeBlend.use(computePassEncoder, targetSrc->bindGroupTexStorage, bindGroupTexStorage, *blendMethod);
dispatchWorkgroups(computePassEncoder);
endRenderPass(computePassEncoder);
};
void WgRenderStorage::compose(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetMsk, WgBindGroupCompositeMethod* composeMethod, WgBindGroupOpacity* opacity)
{
assert(commandEncoder);
assert(targetMsk);
WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder);
mPipelines->computeCompose.use(computePassEncoder, bindGroupTexStorage, targetMsk->bindGroupTexStorage, *composeMethod, *opacity);
dispatchWorkgroups(computePassEncoder);
endRenderPass(computePassEncoder);
};
void WgRenderStorage::dispatchWorkgroups(WGPUComputePassEncoder computePassEncoder)
{
assert(computePassEncoder);
wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, workgroupsCountX, workgroupsCountY, 1);
}
WGPUComputePassEncoder WgRenderStorage::beginComputePass(WGPUCommandEncoder commandEncoder)
{
assert(commandEncoder);
WGPUComputePassDescriptor computePassDesc{};
computePassDesc.nextInChain = nullptr;
computePassDesc.label = "The compute pass composition";
computePassDesc.timestampWrites = nullptr;
return wgpuCommandEncoderBeginComputePass(commandEncoder, &computePassDesc);
};
void WgRenderStorage::endRenderPass(WGPUComputePassEncoder computePassEncoder)
{
assert(computePassEncoder);
wgpuComputePassEncoderEnd(computePassEncoder);
}
//*****************************************************************************
// render storage pool
//*****************************************************************************
WgRenderStorage* WgRenderStoragePool::allocate(WgContext& context, uint32_t w, uint32_t h)
{
WgRenderStorage* renderStorage{};
if (mPool.count > 0) { if (mPool.count > 0) {
renderTarget = mPool.last(); renderStorage = mPool.last();
mPool.pop(); mPool.pop();
} else { } else {
renderTarget = new WgRenderTarget; renderStorage = new WgRenderStorage;
renderTarget->initialize(context, w, h); renderStorage->initialize(context, w, h);
mList.push(renderTarget); mList.push(renderStorage);
} }
return renderTarget; return renderStorage;
}; };
void WgRenderTargetPool::free(WgContext& context, WgRenderTarget* renderTarget) { void WgRenderStoragePool::free(WgContext& context, WgRenderStorage* renderStorage) {
mPool.push(renderTarget); mPool.push(renderStorage);
}; };
void WgRenderTargetPool::release(WgContext& context) void WgRenderStoragePool::release(WgContext& context)
{ {
for (uint32_t i = 0; i < mList.count; i++) { for (uint32_t i = 0; i < mList.count; i++) {
mList[i]->release(context); mList[i]->release(context);

77
src/renderer/wg_engine/tvgWgRenderTarget.h
View file

@ -27,48 +27,65 @@
class WgRenderTarget { class WgRenderTarget {
private: private:
// command buffer // canvas info
WGPURenderPassEncoder mRenderPassEncoder{}; WgBindGroupCanvas mBindGroupCanvas;
// fill and blit data
WgBindGroupCanvas mBindGroupCanvasWnd;
// composition handles
WgMeshData mMeshDataCanvasWnd;
// gpu buffers
WGPUTexture mTextureColor{};
WGPUTexture mTextureStencil{};
WgPipelines* mPipelines{}; // external handle WgPipelines* mPipelines{}; // external handle
public: public:
WGPUSampler sampler{}; WGPUTexture texColor{};
WGPUTextureView textureViewColor{}; WGPUTexture texStencil{};
WGPUTextureView textureViewStencil{}; WGPUTextureView texViewColor{};
WgBindGroupTexture bindGroupTex; WGPUTextureView texViewStencil{};
WgBindGroupStorageTexture bindGroupStorageTex; WgBindGroupTextureStorage bindGroupTexStorage;
WgBindGroupTextureSampled bindGroupTexSampled;
public: public:
void initialize(WgContext& context, uint32_t w, uint32_t h); void initialize(WgContext& context, uint32_t w, uint32_t h);
void release(WgContext& context); void release(WgContext& context);
void beginRenderPass(WGPUCommandEncoder commandEncoder, WGPUTextureView colorAttachement, bool clear); void renderShape(WGPUCommandEncoder commandEncoder, WgRenderDataShape* renderData);
void beginRenderPass(WGPUCommandEncoder commandEncoder, bool clear); void renderPicture(WGPUCommandEncoder commandEncoder, WgRenderDataPicture* renderData);
void endRenderPass(); private:
void drawShape(WGPURenderPassEncoder renderPassEncoder, WgRenderDataShape* renderData);
void drawStroke(WGPURenderPassEncoder renderPassEncoder, WgRenderDataShape* renderData);
void renderShape(WgRenderDataShape* renderData); WGPURenderPassEncoder beginRenderPass(WGPUCommandEncoder commandEncoder);
void renderStroke(WgRenderDataShape* renderData); void endRenderPass(WGPURenderPassEncoder renderPassEncoder);
void renderPicture(WgRenderDataPicture* renderData);
void blit(WgContext& context, WgRenderTarget* renderTargetSrc, WgBindGroupOpacity* bindGroupOpacity);
void blitColor(WgContext& context, WgRenderTarget* renderTargetSrc);
void compose(WgContext& context, WgRenderTarget* renderTargetSrc, WgRenderTarget* renderTargetMsk, CompositeMethod method);
}; };
class WgRenderTargetPool { class WgRenderStorage {
private: private:
Array<WgRenderTarget*> mList; // texture buffers
Array<WgRenderTarget*> mPool; WgPipelines* mPipelines{}; // external handle
public: public:
WgRenderTarget* allocate(WgContext& context, uint32_t w, uint32_t h); WGPUTexture texStorage{};
void free(WgContext& context, WgRenderTarget* renderTarget); WGPUTextureView texViewStorage{};
WgBindGroupTextureStorage bindGroupTexStorage;
uint32_t width{};
uint32_t height{};
uint32_t workgroupsCountX{};
uint32_t workgroupsCountY{};
public:
void initialize(WgContext& context, uint32_t w, uint32_t h);
void release(WgContext& context);
void clear(WGPUCommandEncoder commandEncoder);
void blend(WGPUCommandEncoder commandEncoder, WgRenderTarget* targetSrc, WgBindGroupBlendMethod* blendMethod);
void blend(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc, WgBindGroupBlendMethod* blendMethod);
void compose(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetMsk, WgBindGroupCompositeMethod* composeMethod, WgBindGroupOpacity* opacity);
private:
void dispatchWorkgroups(WGPUComputePassEncoder computePassEncoder);
WGPUComputePassEncoder beginComputePass(WGPUCommandEncoder commandEncoder);
void endRenderPass(WGPUComputePassEncoder computePassEncoder);
};
class WgRenderStoragePool {
private:
Array<WgRenderStorage*> mList;
Array<WgRenderStorage*> mPool;
public:
WgRenderStorage* allocate(WgContext& context, uint32_t w, uint32_t h);
void free(WgContext& context, WgRenderStorage* renderTarget);
void release(WgContext& context); void release(WgContext& context);
}; };

128
src/renderer/wg_engine/tvgWgRenderer.cpp
View file

@ -43,18 +43,22 @@ void WgRenderer::initialize()
{ {
mContext.initialize(); mContext.initialize();
mPipelines.initialize(mContext); mPipelines.initialize(mContext);
mBindGroupOpacityPool.initialize(mContext); mOpacityPool.initialize(mContext);
mBlendMethodPool.initialize(mContext);
mCompositeMethodPool.initialize(mContext);
} }
void WgRenderer::release() void WgRenderer::release()
{ {
mCompositorStack.clear(); mCompositorStack.clear();
mRenderTargetStack.clear(); mRenderStorageStack.clear();
mBindGroupOpacityPool.release(mContext); mRenderStoragePool.release(mContext);
mRenderTargetPool.release(mContext); mCompositeMethodPool.release(mContext);
mRenderTargetRoot.release(mContext); mBlendMethodPool.release(mContext);
mRenderTargetWnd.release(mContext); mOpacityPool.release(mContext);
mRenderStorageRoot.release(mContext);
mRenderTarget.release(mContext);
wgpuSurfaceUnconfigure(mSurface); wgpuSurfaceUnconfigure(mSurface);
wgpuSurfaceRelease(mSurface); wgpuSurfaceRelease(mSurface);
mPipelines.release(); mPipelines.release();
@ -108,7 +112,7 @@ RenderData WgRenderer::prepare(Surface* surface, const RenderMesh* mesh, RenderD
WgShaderTypeBlendSettings blendSettings(surface->cs, opacity); WgShaderTypeBlendSettings blendSettings(surface->cs, opacity);
renderDataPicture->bindGroupPaint.initialize(mContext.device, mContext.queue, modelMat, blendSettings); renderDataPicture->bindGroupPaint.initialize(mContext.device, mContext.queue, modelMat, blendSettings);
} }
// update image data // update image data
if (flags & (RenderUpdateFlag::Path | RenderUpdateFlag::Image)) { if (flags & (RenderUpdateFlag::Path | RenderUpdateFlag::Image)) {
WgGeometryData geometryData; WgGeometryData geometryData;
@ -119,7 +123,7 @@ RenderData WgRenderer::prepare(Surface* surface, const RenderMesh* mesh, RenderD
renderDataPicture->imageData.update(mContext, surface); renderDataPicture->imageData.update(mContext, surface);
renderDataPicture->bindGroupPicture.initialize( renderDataPicture->bindGroupPicture.initialize(
mContext.device, mContext.queue, mContext.device, mContext.queue,
renderDataPicture->imageData.sampler, mContext.samplerLinear,
renderDataPicture->imageData.textureView); renderDataPicture->imageData.textureView);
} }
@ -129,37 +133,41 @@ RenderData WgRenderer::prepare(Surface* surface, const RenderMesh* mesh, RenderD
bool WgRenderer::preRender() bool WgRenderer::preRender()
{ {
// command encoder descriptor
WGPUCommandEncoderDescriptor commandEncoderDesc{}; WGPUCommandEncoderDescriptor commandEncoderDesc{};
commandEncoderDesc.nextInChain = nullptr; commandEncoderDesc.nextInChain = nullptr;
commandEncoderDesc.label = "The command encoder"; commandEncoderDesc.label = "The command encoder";
mCommandEncoder = wgpuDeviceCreateCommandEncoder(mContext.device, &commandEncoderDesc); mCommandEncoder = wgpuDeviceCreateCommandEncoder(mContext.device, &commandEncoderDesc);
// render datas mRenderStorageRoot.clear(mCommandEncoder);
mRenderTargetStack.push(&mRenderTargetRoot); mRenderStorageStack.push(&mRenderStorageRoot);
mRenderTargetRoot.beginRenderPass(mCommandEncoder, true);
return true; return true;
} }
bool WgRenderer::renderShape(RenderData data) bool WgRenderer::renderShape(RenderData data)
{ {
mRenderTargetStack.last()->renderShape((WgRenderDataShape *)data); // render shape to render target
mRenderTargetStack.last()->renderStroke((WgRenderDataShape *)data); mRenderTarget.renderShape(mCommandEncoder, (WgRenderDataShape *)data);
// blend shape with current render storage
WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod);
mRenderStorageStack.last()->blend(mCommandEncoder, &mRenderTarget, blendMethod);
return true; return true;
} }
bool WgRenderer::renderImage(RenderData data) bool WgRenderer::renderImage(RenderData data)
{ {
mRenderTargetStack.last()->renderPicture((WgRenderDataPicture *)data); // render image to render target
mRenderTarget.renderPicture(mCommandEncoder, (WgRenderDataPicture *)data);
// blend image with current render storage
WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod);
mRenderStorageStack.last()->blend(mCommandEncoder, &mRenderTarget, blendMethod);
return true; return true;
} }
bool WgRenderer::postRender() bool WgRenderer::postRender()
{ {
mRenderTargetRoot.endRenderPass(); mRenderStorageStack.pop();
mRenderTargetStack.pop();
mContext.executeCommandEncoder(mCommandEncoder); mContext.executeCommandEncoder(mCommandEncoder);
wgpuCommandEncoderRelease(mCommandEncoder); wgpuCommandEncoderRelease(mCommandEncoder);
return true; return true;
@ -193,6 +201,7 @@ bool WgRenderer::viewport(TVG_UNUSED const RenderRegion& vp)
bool WgRenderer::blend(TVG_UNUSED BlendMethod method) bool WgRenderer::blend(TVG_UNUSED BlendMethod method)
{ {
mBlendMethod = method;
return false; return false;
} }
@ -213,17 +222,25 @@ bool WgRenderer::sync()
{ {
WGPUSurfaceTexture backBuffer{}; WGPUSurfaceTexture backBuffer{};
wgpuSurfaceGetCurrentTexture(mSurface, &backBuffer); wgpuSurfaceGetCurrentTexture(mSurface, &backBuffer);
WGPUTextureView backBufferView = mContext.createTextureView2d(backBuffer.texture, "Surface texture view");
WGPUCommandEncoderDescriptor commandEncoderDesc{}; WGPUCommandEncoderDescriptor commandEncoderDesc{};
commandEncoderDesc.nextInChain = nullptr; commandEncoderDesc.nextInChain = nullptr;
commandEncoderDesc.label = "The command encoder"; commandEncoderDesc.label = "The command encoder";
WGPUCommandEncoder commandEncoder = wgpuDeviceCreateCommandEncoder(mContext.device, &commandEncoderDesc); WGPUCommandEncoder commandEncoder = wgpuDeviceCreateCommandEncoder(mContext.device, &commandEncoderDesc);
mRenderTargetWnd.beginRenderPass(commandEncoder, backBufferView, true);
mRenderTargetWnd.blitColor(mContext, &mRenderTargetRoot); WGPUImageCopyTexture source{};
mRenderTargetWnd.endRenderPass(); source.texture = mRenderStorageRoot.texStorage;
WGPUImageCopyTexture dest{};
dest.texture = backBuffer.texture;
WGPUExtent3D copySize{};
copySize.width = mTargetSurface.w;
copySize.height = mTargetSurface.h;
copySize.depthOrArrayLayers = 1;
wgpuCommandEncoderCopyTextureToTexture(commandEncoder, &source, &dest, &copySize);
mContext.executeCommandEncoder(commandEncoder); mContext.executeCommandEncoder(commandEncoder);
wgpuCommandEncoderRelease(commandEncoder); wgpuCommandEncoderRelease(commandEncoder);
wgpuTextureViewRelease(backBufferView);
wgpuSurfacePresent(mSurface); wgpuSurfacePresent(mSurface);
return true; return true;
} }
@ -236,7 +253,7 @@ bool WgRenderer::target(uint32_t* buffer, uint32_t stride, uint32_t w, uint32_t
mTargetSurface.w = w; mTargetSurface.w = w;
mTargetSurface.h = h; mTargetSurface.h = h;
mRenderTargetRoot.initialize(mContext, w, h); mRenderTarget.initialize(mContext, w, h);
return true; return true;
} }
@ -266,7 +283,7 @@ bool WgRenderer::target(void* window, uint32_t w, uint32_t h)
surfaceConfiguration.nextInChain = nullptr; surfaceConfiguration.nextInChain = nullptr;
surfaceConfiguration.device = mContext.device; surfaceConfiguration.device = mContext.device;
surfaceConfiguration.format = WGPUTextureFormat_RGBA8Unorm; surfaceConfiguration.format = WGPUTextureFormat_RGBA8Unorm;
surfaceConfiguration.usage = WGPUTextureUsage_RenderAttachment; surfaceConfiguration.usage = WGPUTextureUsage_RenderAttachment | WGPUTextureUsage_CopyDst;
surfaceConfiguration.viewFormatCount = 0; surfaceConfiguration.viewFormatCount = 0;
surfaceConfiguration.viewFormats = nullptr; surfaceConfiguration.viewFormats = nullptr;
surfaceConfiguration.alphaMode = WGPUCompositeAlphaMode_Auto; surfaceConfiguration.alphaMode = WGPUCompositeAlphaMode_Auto;
@ -275,8 +292,8 @@ bool WgRenderer::target(void* window, uint32_t w, uint32_t h)
surfaceConfiguration.presentMode = WGPUPresentMode_Mailbox; surfaceConfiguration.presentMode = WGPUPresentMode_Mailbox;
wgpuSurfaceConfigure(mSurface, &surfaceConfiguration); wgpuSurfaceConfigure(mSurface, &surfaceConfiguration);
mRenderTargetWnd.initialize(mContext, w, h); mRenderTarget.initialize(mContext, w, h);
mRenderTargetRoot.initialize(mContext, w, h); mRenderStorageRoot.initialize(mContext, w, h);
return true; return true;
} }
@ -291,17 +308,14 @@ Compositor* WgRenderer::target(TVG_UNUSED const RenderRegion& region, TVG_UNUSED
bool WgRenderer::beginComposite(TVG_UNUSED Compositor* cmp, TVG_UNUSED CompositeMethod method, TVG_UNUSED uint8_t opacity) bool WgRenderer::beginComposite(TVG_UNUSED Compositor* cmp, TVG_UNUSED CompositeMethod method, TVG_UNUSED uint8_t opacity)
{ {
// save current composition settings // save current composition settings
cmp->method = method; cmp->method = method;
cmp->opacity = opacity; cmp->opacity = opacity;
// end current render target // allocate new render storage and push it to top of render tree
mRenderTargetStack.last()->endRenderPass(); WgRenderStorage* renderStorage = mRenderStoragePool.allocate(mContext, mTargetSurface.w, mTargetSurface.h);
renderStorage->clear(mCommandEncoder);
// create new render target and begin new render pass mRenderStorageStack.push(renderStorage);
WgRenderTarget* renderTarget = mRenderTargetPool.allocate(mContext, mTargetSurface.w, mTargetSurface.h);
renderTarget->beginRenderPass(mCommandEncoder, true);
mRenderTargetStack.push(renderTarget);
return true; return true;
} }
@ -310,41 +324,35 @@ bool WgRenderer::beginComposite(TVG_UNUSED Compositor* cmp, TVG_UNUSED Composite
bool WgRenderer::endComposite(TVG_UNUSED Compositor* cmp) bool WgRenderer::endComposite(TVG_UNUSED Compositor* cmp)
{ {
if (cmp->method == CompositeMethod::None) { if (cmp->method == CompositeMethod::None) {
// end current render pass
mRenderTargetStack.last()->endRenderPass();
// get two last render targets // get two last render targets
WgRenderTarget* renderTargetSrc = mRenderTargetStack.last(); WgRenderStorage* renderStorageSrc = mRenderStorageStack.last();
mRenderTargetStack.pop(); mRenderStorageStack.pop();
// apply current render target // blent scene to current render storage
WgRenderTarget* renderTarget = mRenderTargetStack.last(); WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod);
renderTarget->beginRenderPass(mCommandEncoder, false); mRenderStorageStack.last()->blend(mCommandEncoder, renderStorageSrc, blendMethod);
// blit scene to current render tartget
WgBindGroupOpacity* mBindGroupOpacity = mBindGroupOpacityPool.allocate(mContext, cmp->opacity);
renderTarget->blit(mContext, renderTargetSrc, mBindGroupOpacity);
// back render targets to the pool // back render targets to the pool
mRenderTargetPool.free(mContext, renderTargetSrc); mRenderStoragePool.free(mContext, renderStorageSrc);
} else { } else {
// end current render pass
mRenderTargetStack.last()->endRenderPass();
// get two last render targets // get two last render targets
WgRenderTarget* renderTargetSrc = mRenderTargetStack.last(); WgRenderStorage* renderStorageSrc = mRenderStorageStack.last();
mRenderTargetStack.pop(); mRenderStorageStack.pop();
WgRenderTarget* renderTargetMsk = mRenderTargetStack.last(); WgRenderStorage* renderStorageMsk = mRenderStorageStack.last();
mRenderTargetStack.pop(); mRenderStorageStack.pop();
// apply current render target // compose shape and mask
WgRenderTarget* renderTarget = mRenderTargetStack.last(); WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, cmp->opacity);
renderTarget->beginRenderPass(mCommandEncoder, false); WgBindGroupCompositeMethod* composeMethod = mCompositeMethodPool.allocate(mContext, cmp->method);
renderTarget->compose(mContext, renderTargetSrc, renderTargetMsk, cmp->method); renderStorageSrc->compose(mCommandEncoder, renderStorageMsk, composeMethod, opacity);
// blent scene to current render storage
WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod);
mRenderStorageStack.last()->blend(mCommandEncoder, renderStorageSrc, blendMethod);
// back render targets to the pool // back render targets to the pool
mRenderTargetPool.free(mContext, renderTargetSrc); mRenderStoragePool.free(mContext, renderStorageSrc);
mRenderTargetPool.free(mContext, renderTargetMsk); mRenderStoragePool.free(mContext, renderStorageMsk);
} }
// delete current compositor // delete current compositor

21
src/renderer/wg_engine/tvgWgRenderer.h
View file

@ -62,19 +62,24 @@ public:
private: private:
// render handles // render handles
WGPUCommandEncoder mCommandEncoder{}; WGPUCommandEncoder mCommandEncoder{};
Array<Compositor*> mCompositorStack; WgRenderTarget mRenderTarget;
Array<WgRenderTarget*> mRenderTargetStack; WgRenderStorage mRenderStorageRoot;
WgRenderStoragePool mRenderStoragePool;
WgBindGroupOpacityPool mOpacityPool;
WgBindGroupBlendMethodPool mBlendMethodPool;
WgBindGroupCompositeMethodPool mCompositeMethodPool;
// render object pools // render tree stacks
WgRenderTargetPool mRenderTargetPool; Array<Compositor*> mCompositorStack;
WgBindGroupOpacityPool mBindGroupOpacityPool; Array<WgRenderStorage*> mRenderStorageStack;
// native window handles
WGPUSurface mSurface{};
private: private:
WgContext mContext; WgContext mContext;
WgPipelines mPipelines; WgPipelines mPipelines;
WgRenderTarget mRenderTargetRoot;
WgRenderTarget mRenderTargetWnd;
WGPUSurface mSurface{};
Surface mTargetSurface; Surface mTargetSurface;
BlendMethod mBlendMethod{};
}; };
#endif /* _TVG_WG_RENDERER_H_ */ #endif /* _TVG_WG_RENDERER_H_ */

94
src/renderer/wg_engine/tvgWgShaderSrc.cpp
View file

@ -712,18 +712,92 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f {
// cShaderSource_PipelineComputeBlend // cShaderSource_PipelineComputeBlend
//************************************************************************ //************************************************************************
// pipeline shader modules blend (simple example) // pipeline shader modules clear
const char* cShaderSource_PipelineComputeBlend = R"( const char* cShaderSource_PipelineComputeClear = R"(
@group(0) @binding(0) var imageSrc : texture_storage_2d<rgba8unorm, write>; @group(0) @binding(0) var imageDst : texture_storage_2d<rgba8unorm, read_write>;
@group(1) @binding(0) var imageDst : texture_storage_2d<rgba8unorm, write>;
@compute @workgroup_size(8, 8) @compute @workgroup_size(8, 8)
fn cs_main( @builtin(global_invocation_id) id: vec3u) { fn cs_main( @builtin(global_invocation_id) id: vec3u) {
let texSize = textureDimensions(imageSrc); textureStore(imageDst, id.xy, vec4f(0.0, 0.0, 0.0, 0.0));
if ((id.x >= texSize.x) || (id.y >= texSize.y)) { return; }; }
)";
// var src = textureLoad(imageSrc, id.xy);
textureStore(imageDst, id.xy, vec4(0.5, 1.0, 0.1, 1.0));
// pipeline shader modules blend
const char* cShaderSource_PipelineComputeBlend = R"(
@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(2) @binding(0) var<uniform> blendMethod : u32;
@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 srcColor = textureLoad(imageSrc, id.xy);
if (srcColor.a == 0.0) { return; };
let dstColor = textureLoad(imageDst, id.xy);
let Sa: f32 = srcColor.a;
let Da: f32 = dstColor.a;
let S: vec3f = srcColor.xyz;
let D: vec3f = dstColor.xyz;
let One: vec3f = vec3(1.0);
var color: vec3f = vec3f(0.0, 0.0, 0.0);
switch blendMethod {
/* Normal */ case 0u: { color = (Sa * S) + (1.0 - Sa) * D; }
/* Add */ case 1u: { color = (S + D); }
/* Screen */ case 2u: { color = (S + D) - (S * D); }
/* Multiply */ case 3u: { color = (S * D); }
/* Overlay */ case 4u: { color = S * D; }
/* Difference */ case 5u: { color = abs(S - D); }
/* Exclusion */ case 6u: { color = S + D - (2 * S * D); }
/* SrcOver */ case 7u: { color = S; }
/* Darken */ case 8u: { color = min(S, D); }
/* Lighten */ case 9u: { color = max(S, D); }
/* ColorDodge */ case 10u: { color = D / (One - S); }
/* ColorBurn */ case 11u: { color = One - (One - D) / S; }
/* HardLight */ case 12u: { color = (Sa * Da) - 2.0 * (Da - S) * (Sa - D); }
/* SoftLight */ case 13u: { color = (One - 2 * S) * (D * D) + (2 * S * D); }
default: { color = (Sa * S) + (1.0 - Sa) * D; }
}
textureStore(imageDst, id.xy, vec4f(color, Sa));
}
)";
// pipeline shader modules compose
const char* cShaderSource_PipelineComputeCompose = R"(
@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(2) @binding(0) var<uniform> composeMethod : u32;
@group(3) @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 colorSrc = textureLoad(imageSrc, id.xy);
let colorMsk = textureLoad(imageMsk, id.xy);
var color: vec3f = colorSrc.xyz;
var alpha: f32 = colorMsk.a;
switch composeMethod {
/* None */ case 0u: { color = colorSrc.xyz; }
/* ClipPath */ case 1u: { if (colorMsk.a == 0) { alpha = 0.0; }; }
/* AlphaMask */ case 2u: { color = mix(colorMsk.xyz, colorSrc.xyz, colorSrc.a * colorMsk.b); }
/* InvAlphaMask */ case 3u: { color = mix(colorSrc.xyz, colorMsk.xyz, colorSrc.a * colorMsk.b); alpha = 1.0 - colorMsk.b; }
/* LumaMask */ case 4u: { color = colorSrc.xyz * (0.299 * colorMsk.r + 0.587 * colorMsk.g + 0.114 * colorMsk.b); }
/* InvLumaMask */ case 5u: { color = colorSrc.xyz * (1.0 - (0.299 * colorMsk.r + 0.587 * colorMsk.g + 0.114 * colorMsk.b)); alpha = 1.0 - colorMsk.b; }
/* AddMask */ case 6u: { color = colorSrc.xyz * colorSrc.a + colorMsk.xyz * (1.0 - colorSrc.a); }
/* SubtractMask */ case 7u: { color = colorSrc.xyz * colorSrc.a - colorMsk.xyz * (1.0 - colorSrc.a); }
/* IntersectMask */ case 8u: { color = colorSrc.xyz * min(colorSrc.a, colorMsk.a); }
/* DifferenceMask */ case 9u: { color = abs(colorSrc.xyz - colorMsk.xyz * (1.0 - colorMsk.a)); }
default: { color = colorSrc.xyz; }
}
textureStore(imageSrc, id.xy, vec4f(color, alpha * opacity));
} }
)"; )";

4
src/renderer/wg_engine/tvgWgShaderSrc.h
View file

@ -54,7 +54,9 @@ extern const char* cShaderSource_PipelineCompDifferenceMask;
// compute shader modules // compute shader modules
//***************************************************************************** //*****************************************************************************
// pipeline shader modules blend // pipeline shader modules clear, compose and blend
extern const char* cShaderSource_PipelineComputeClear;
extern const char* cShaderSource_PipelineComputeBlend; extern const char* cShaderSource_PipelineComputeBlend;
extern const char* cShaderSource_PipelineComputeCompose;
#endif // _TVG_WG_SHADER_SRC_H_ #endif // _TVG_WG_SHADER_SRC_H_