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

wg_engine: composition and blend optimization

Browse source 
* bind groups creation in real time removed - performance boost
* blend and composition shaders decomposed - performance boost
* shader modules and pipeline layouts generalized - less memory usage
* shared single stencil buffer used - less memory usage
* bind groups usage simplified
* general context API simplified and generalized
* all rendering logic moved into new composition class
* ready for hardware MSAA (in next steps)
* ready for direct mask applience (in next steps)
This commit is contained in:
Sergii Liebodkin 2024-08-06 19:47:49 +00:00 • committed by Hermet Park
parent 997093be96
commit f281cc9e92
21 changed files with 1949 additions and 3716 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

2
src/renderer/wg_engine/meson.build Normal file → Executable file
View file

@ -1,6 +1,7 @@
source_file = [
'tvgWgBindGroups.h',
'tvgWgCommon.h',
'tvgWgCompositor.h',
'tvgWgGeometry.h',
'tvgWgPipelines.h',
'tvgWgRenderData.h',
@ -10,6 +11,7 @@ source_file = [
'tvgWgShaderTypes.h',
'tvgWgBindGroups.cpp',
'tvgWgCommon.cpp',
'tvgWgCompositor.cpp',
'tvgWgGeometry.cpp',
'tvgWgPipelines.cpp',
'tvgWgRenderData.cpp',

1047
src/renderer/wg_engine/tvgWgBindGroups.cpp Normal file → Executable file
View file

File diff suppressed because it is too large Load diff

312
src/renderer/wg_engine/tvgWgBindGroups.h Normal file → Executable file
View file

@ -24,300 +24,34 @@
#define _TVG_WG_BIND_GROUPS_H_
#include "tvgWgCommon.h"
#include "tvgWgShaderTypes.h"
// @group(0)
struct WgBindGroupCanvas : public WgBindGroup
{
static WGPUBindGroupLayout layout;
static WGPUBindGroupLayout getLayout(WGPUDevice device);
static void releaseLayout();
WGPUBuffer uBufferViewMat{};
void initialize(WGPUDevice device, WGPUQueue queue,
WgShaderTypeMat4x4f& uViewMat);
void release();
};
// @group(1)
struct WgBindGroupPaint : public WgBindGroup
{
static WGPUBindGroupLayout layout;
static WGPUBindGroupLayout getLayout(WGPUDevice device);
static void releaseLayout();
WGPUBuffer uBufferModelMat{};
WGPUBuffer uBufferBlendSettings{};
void initialize(WGPUDevice device, WGPUQueue queue,
WgShaderTypeMat4x4f& uModelMat,
WgShaderTypeBlendSettings& uBlendSettings);
void release();
};
// @group(2)
struct WgBindGroupSolidColor : public WgBindGroup
{
static WGPUBindGroupLayout layout;
static WGPUBindGroupLayout getLayout(WGPUDevice device);
static void releaseLayout();
WGPUBuffer uBufferSolidColor{};
void initialize(WGPUDevice device, WGPUQueue queue,
WgShaderTypeSolidColor &uSolidColor);
void release();
};
// @group(2)
struct WgBindGroupLinearGradient : public WgBindGroup
{
static WGPUBindGroupLayout layout;
static WGPUBindGroupLayout getLayout(WGPUDevice device);
static void releaseLayout();
WGPUBuffer uBufferLinearGradient{};
void initialize(WGPUDevice device, WGPUQueue queue,
WgShaderTypeLinearGradient &uLinearGradient);
void release();
};
// @group(2)
struct WgBindGroupRadialGradient : public WgBindGroup
{
static WGPUBindGroupLayout layout;
static WGPUBindGroupLayout getLayout(WGPUDevice device);
static void releaseLayout();
WGPUBuffer uBufferRadialGradient{};
void initialize(WGPUDevice device, WGPUQueue queue,
WgShaderTypeRadialGradient &uRadialGradient);
void release();
};
// @group(2)
struct WgBindGroupPicture : public WgBindGroup
{
static WGPUBindGroupLayout layout;
static WGPUBindGroupLayout getLayout(WGPUDevice device);
static void releaseLayout();
void initialize(WGPUDevice device, WGPUQueue queue,
WGPUSampler uSampler,
WGPUTextureView uTextureView);
void release();
};
// @group(0 or 1)
struct WgBindGroupTexture : public WgBindGroup
{
static WGPUBindGroupLayout layout;
static WGPUBindGroupLayout getLayout(WGPUDevice device);
static void releaseLayout();
void initialize(WGPUDevice device, WGPUQueue queue,
WGPUTextureView uTexture);
void release();
};
// @group(0 or 1)
struct WgBindGroupTextureStorageRgbaWO : public WgBindGroup
{
static WGPUBindGroupLayout layout;
static WGPUBindGroupLayout getLayout(WGPUDevice device);
static void releaseLayout();
void initialize(WGPUDevice device, WGPUQueue queue,
WGPUTextureView uTexture);
void release();
};
// @group(0 or 1)
struct WgBindGroupTextureStorageRgbaRO : public WgBindGroup
{
static WGPUBindGroupLayout layout;
static WGPUBindGroupLayout getLayout(WGPUDevice device);
static void releaseLayout();
void initialize(WGPUDevice device, WGPUQueue queue,
WGPUTextureView uTexture);
void release();
};
// @group(0 or 1)
struct WgBindGroupTextureStorageBgraRO : public WgBindGroup
{
static WGPUBindGroupLayout layout;
static WGPUBindGroupLayout getLayout(WGPUDevice device);
static WGPUBindGroupLayout getLayoutRO(WGPUDevice device);
static WGPUBindGroupLayout getLayoutWO(WGPUDevice device);
static void releaseLayout();
void initialize(WGPUDevice device, WGPUQueue queue,
WGPUTextureView uTexture);
void release();
};
// @group(0 or 1)
struct WgBindGroupTextureStorageBgraWO : public WgBindGroup
{
static WGPUBindGroupLayout layout;
static WGPUBindGroupLayout getLayout(WGPUDevice device);
static WGPUBindGroupLayout getLayoutRO(WGPUDevice device);
static WGPUBindGroupLayout getLayoutWO(WGPUDevice device);
static void releaseLayout();
void initialize(WGPUDevice device, WGPUQueue queue,
WGPUTextureView uTexture);
void release();
};
// @group(0 or 1)
struct WgBindGroupTextureSampled : public WgBindGroup
{
static WGPUBindGroupLayout layout;
static WGPUBindGroupLayout getLayout(WGPUDevice device);
static void releaseLayout();
void initialize(WGPUDevice device, WGPUQueue queue,
WGPUSampler uSampler,
WGPUTextureView uTexture);
void release();
};
// @group(0)
struct WgBindGroupTexBlend : public WgBindGroup
{
static WGPUBindGroupLayout layout;
static WGPUBindGroupLayout getLayout(WGPUDevice device);
static void releaseLayout();
void initialize(WGPUDevice device, WGPUQueue queue,
WGPUTextureView uTexSrc,
WGPUTextureView uTexDst,
WGPUTextureView uTexTrg);
void release();
};
// @group(0)
struct WgBindGroupTexBlendMask : public WgBindGroup
{
static WGPUBindGroupLayout layout;
static WGPUBindGroupLayout getLayout(WGPUDevice device);
static void releaseLayout();
void initialize(WGPUDevice device, WGPUQueue queue,
WGPUTextureView uTexSrc,
WGPUTextureView uTexMsk,
WGPUTextureView uTexDst,
WGPUTextureView uTexTrg);
void release();
};
// @group(0)
struct WgBindGroupTexMaskCompose : public WgBindGroup
{
static WGPUBindGroupLayout layout;
static WGPUBindGroupLayout getLayout(WGPUDevice device);
static void releaseLayout();
void initialize(WGPUDevice device, WGPUQueue queue,
WGPUTextureView uTexMsk0,
WGPUTextureView uTexMsk1,
WGPUTextureView uTexTrg);
void release();
};
// @group(0)
struct WgBindGroupTexCompose : public WgBindGroup
{
static WGPUBindGroupLayout layout;
static WGPUBindGroupLayout getLayout(WGPUDevice device);
static void releaseLayout();
void initialize(WGPUDevice device, WGPUQueue queue,
WGPUTextureView uTexSrc,
WGPUTextureView uTexMsk,
WGPUTextureView uTexDst,
WGPUTextureView uTexTrg);
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
//************************************************************************
class WgBindGroupOpacityPool
{
class WgBindGroupLayouts {
private:
WgBindGroupOpacity* mPool[256]{};
WGPUDevice device{};
public:
WGPUBindGroupLayout layoutTexSampled{};
WGPUBindGroupLayout layoutTexScreen1WO{};
WGPUBindGroupLayout layoutTexStrorage1WO{};
WGPUBindGroupLayout layoutTexStrorage1RO{};
WGPUBindGroupLayout layoutTexStrorage2RO{};
WGPUBindGroupLayout layoutTexStrorage3RO{};
WGPUBindGroupLayout layoutBuffer1Un{};
WGPUBindGroupLayout layoutBuffer2Un{};
WGPUBindGroupLayout layoutBuffer3Un{};
public:
WGPUBindGroup createBindGroupTexSampled(WGPUSampler sampler, WGPUTextureView texView);
WGPUBindGroup createBindGroupScreen1WO(WGPUTextureView texView);
WGPUBindGroup createBindGroupStrorage1WO(WGPUTextureView texView);
WGPUBindGroup createBindGroupStrorage1RO(WGPUTextureView texView);
WGPUBindGroup createBindGroupStrorage2RO(WGPUTextureView texView0, WGPUTextureView texView1);
WGPUBindGroup createBindGroupStrorage3RO(WGPUTextureView texView0, WGPUTextureView texView1, WGPUTextureView texView2);
WGPUBindGroup createBindGroupBuffer1Un(WGPUBuffer buff);
WGPUBindGroup createBindGroupBuffer2Un(WGPUBuffer buff0, WGPUBuffer buff1);
WGPUBindGroup createBindGroupBuffer3Un(WGPUBuffer buff0, WGPUBuffer buff1, WGPUBuffer buff2);
void releaseBindGroup(WGPUBindGroup& bindGroup);
public:
void initialize(WgContext& context);
void release(WgContext& context);
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_

825
src/renderer/wg_engine/tvgWgCommon.cpp Normal file → Executable file
View file

@ -20,16 +20,13 @@
* SOFTWARE.
*/
#include "tvgWgCommon.h"
#ifdef __EMSCRIPTEN__
#include <emscripten.h>
#endif
#include "tvgWgCommon.h"
#include "tvgArray.h"
#include <iostream>
//*****************************************************************************
// context
//*****************************************************************************
void WgContext::initialize(WGPUInstance instance, WGPUSurface surface)
{
assert(instance);
@ -39,198 +36,119 @@ void WgContext::initialize(WGPUInstance instance, WGPUSurface surface)
this->instance = instance;
this->surface = surface;
// request adapter options
WGPURequestAdapterOptions requestAdapterOptions{};
requestAdapterOptions.nextInChain = nullptr;
requestAdapterOptions.compatibleSurface = surface;
requestAdapterOptions.powerPreference = WGPUPowerPreference_HighPerformance;
requestAdapterOptions.forceFallbackAdapter = false;
// on adapter request ended function
auto onAdapterRequestEnded = [](WGPURequestAdapterStatus status, WGPUAdapter adapter, char const * message, void * pUserData) {
if (status != WGPURequestAdapterStatus_Success)
TVGERR("WG_RENDERER", "Adapter request: %s", message);
*((WGPUAdapter*)pUserData) = adapter;
};
// request adapter
const WGPURequestAdapterOptions requestAdapterOptions { .nextInChain = nullptr, .compatibleSurface = surface, .powerPreference = WGPUPowerPreference_HighPerformance, .forceFallbackAdapter = false };
auto onAdapterRequestEnded = [](WGPURequestAdapterStatus status, WGPUAdapter adapter, char const * message, void * pUserData) { *((WGPUAdapter*)pUserData) = adapter; };
wgpuInstanceRequestAdapter(instance, &requestAdapterOptions, onAdapterRequestEnded, &adapter);
#ifdef __EMSCRIPTEN__
while (!adapter) emscripten_sleep(10);
#endif
assert(adapter);
// get feature names
// get adapter and surface properties
WGPUFeatureName featureNames[32]{};
size_t featuresCount = wgpuAdapterEnumerateFeatures(adapter, featureNames);
preferredFormat = wgpuSurfaceGetPreferredFormat(surface, adapter);
// request device
WGPUDeviceDescriptor deviceDesc{};
deviceDesc.nextInChain = nullptr;
deviceDesc.label = "The device";
deviceDesc.requiredFeatureCount = featuresCount;
deviceDesc.requiredFeatures = featureNames;
deviceDesc.requiredLimits = nullptr;
deviceDesc.defaultQueue.nextInChain = nullptr;
deviceDesc.defaultQueue.label = "The default queue";
deviceDesc.deviceLostCallback = nullptr;
deviceDesc.deviceLostUserdata = nullptr;
// on device request ended function
auto onDeviceRequestEnded = [](WGPURequestDeviceStatus status, WGPUDevice device, char const * message, void * pUserData) {
if (status != WGPURequestDeviceStatus_Success)
TVGERR("WG_RENDERER", "Device request: %s", message);
*((WGPUDevice*)pUserData) = device;
};
// request device
const WGPUDeviceDescriptor deviceDesc { .nextInChain = nullptr, .label = "The device", .requiredFeatureCount = featuresCount, .requiredFeatures = featureNames };
auto onDeviceRequestEnded = [](WGPURequestDeviceStatus status, WGPUDevice device, char const * message, void * pUserData) { *((WGPUDevice*)pUserData) = device; };
wgpuAdapterRequestDevice(adapter, &deviceDesc, onDeviceRequestEnded, &device);
#ifdef __EMSCRIPTEN__
while (!device) emscripten_sleep(10);
#endif
assert(device);
// on device error function
auto onDeviceError = [](WGPUErrorType type, char const* message, void* pUserData) {
TVGERR("WG_RENDERER", "Uncaptured device error: %s", message);
// TODO: remove direct error message
std::cout << message << std::endl;
};
// set device error handling
// device uncaptured error callback
auto onDeviceError = [](WGPUErrorType type, char const* message, void* pUserData) { std::cout << message << std::endl; };
wgpuDeviceSetUncapturedErrorCallback(device, onDeviceError, nullptr);
// get current queue
queue = wgpuDeviceGetQueue(device);
assert(queue);
// create default nearest and linear samplers
// create shared webgpu assets
allocateBufferIndexFan(32768);
samplerNearest = createSampler(WGPUFilterMode_Nearest, WGPUMipmapFilterMode_Nearest);
assert(samplerNearest);
samplerLinear = createSampler(WGPUFilterMode_Linear, WGPUMipmapFilterMode_Linear);
assert(samplerNearest);
assert(samplerLinear);
allocateIndexBufferFan(1024);
assert(indexBufferFan);
}
void WgContext::release()
{
releaseBuffer(indexBufferFan);
releaseSampler(samplerNearest);
releaseSampler(samplerLinear);
if (device) {
wgpuDeviceDestroy(device);
wgpuDeviceRelease(device);
device = nullptr;
}
if (adapter) {
wgpuAdapterRelease(adapter);
adapter = nullptr;
}
if (surface) {
wgpuSurfaceUnconfigure(surface);
wgpuSurfaceRelease(surface);
surface = nullptr;
}
if (instance) {
wgpuInstanceRelease(instance);
instance = nullptr;
}
releaseBuffer(bufferIndexFan);
wgpuDeviceRelease(device);
wgpuAdapterRelease(adapter);
}
void WgContext::executeCommandEncoder(WGPUCommandEncoder commandEncoder)
WGPUSampler WgContext::createSampler(WGPUFilterMode filter, WGPUMipmapFilterMode mipmapFilter)
{
// command buffer descriptor
WGPUCommandBufferDescriptor commandBufferDesc{};
commandBufferDesc.nextInChain = nullptr;
commandBufferDesc.label = "The command buffer";
WGPUCommandBuffer commandsBuffer = nullptr;
commandsBuffer = wgpuCommandEncoderFinish(commandEncoder, &commandBufferDesc);
wgpuQueueSubmit(queue, 1, &commandsBuffer);
wgpuCommandBufferRelease(commandsBuffer);
}
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;
const WGPUSamplerDescriptor samplerDesc {
.magFilter = filter, .minFilter = filter, .mipmapFilter = mipmapFilter,
.lodMinClamp = 0.0f, .lodMaxClamp = 32.0f, .maxAnisotropy = 1
};
return wgpuDeviceCreateSampler(device, &samplerDesc);
}
WGPUTexture WgContext::createTexture2d(WGPUTextureUsageFlags usage, WGPUTextureFormat format, uint32_t width, uint32_t height, char const * label) {
WGPUTextureDescriptor textureDesc{};
textureDesc.nextInChain = nullptr;
textureDesc.label = label;
textureDesc.usage = usage;
textureDesc.dimension = WGPUTextureDimension_2D;
textureDesc.size = { width, height, 1 };
textureDesc.format = format;
textureDesc.mipLevelCount = 1;
textureDesc.sampleCount = 1;
textureDesc.viewFormatCount = 0;
textureDesc.viewFormats = nullptr;
WGPUTexture WgContext::createTexture(uint32_t width, uint32_t height, WGPUTextureFormat format)
{
const WGPUTextureDescriptor textureDesc {
.usage = WGPUTextureUsage_CopyDst | WGPUTextureUsage_TextureBinding,
.dimension = WGPUTextureDimension_2D, .size = { width, height, 1 },
.format = format, .mipLevelCount = 1, .sampleCount = 1
};
return wgpuDeviceCreateTexture(device, &textureDesc);
}
WGPUTexture WgContext::createTexture2dMS(WGPUTextureUsageFlags usage, WGPUTextureFormat format, uint32_t width, uint32_t height, uint32_t sc, char const * label)
WGPUTexture WgContext::createTexStorage(uint32_t width, uint32_t height, WGPUTextureFormat format, uint32_t sc)
{
WGPUTextureDescriptor textureDesc{};
textureDesc.nextInChain = nullptr;
textureDesc.label = label;
textureDesc.usage = usage;
textureDesc.dimension = WGPUTextureDimension_2D;
textureDesc.size = { width, height, 1 };
textureDesc.format = format;
textureDesc.mipLevelCount = 1;
textureDesc.sampleCount = sc;
textureDesc.viewFormatCount = 0;
textureDesc.viewFormats = nullptr;
const WGPUTextureDescriptor textureDesc {
.usage = WGPUTextureUsage_CopySrc | WGPUTextureUsage_CopyDst | WGPUTextureUsage_TextureBinding | WGPUTextureUsage_StorageBinding | WGPUTextureUsage_RenderAttachment,
.dimension = WGPUTextureDimension_2D, .size = { width, height, 1 },
.format = format, .mipLevelCount = 1, .sampleCount = sc
};
return wgpuDeviceCreateTexture(device, &textureDesc);
}
WGPUTextureView WgContext::createTextureView2d(WGPUTexture texture, char const * label)
WGPUTexture WgContext::createTexStencil(uint32_t width, uint32_t height, WGPUTextureFormat format, uint32_t sc)
{
WGPUTextureViewDescriptor textureViewDescColor{};
textureViewDescColor.nextInChain = nullptr;
textureViewDescColor.label = label;
textureViewDescColor.format = wgpuTextureGetFormat(texture);
textureViewDescColor.dimension = WGPUTextureViewDimension_2D;
textureViewDescColor.baseMipLevel = 0;
textureViewDescColor.mipLevelCount = 1;
textureViewDescColor.baseArrayLayer = 0;
textureViewDescColor.arrayLayerCount = 1;
textureViewDescColor.aspect = WGPUTextureAspect_All;
return wgpuTextureCreateView(texture, &textureViewDescColor);
};
WGPUBuffer WgContext::createBuffer(WGPUBufferUsageFlags usage, uint64_t size,char const * label)
{
WGPUBufferDescriptor bufferDesc{};
bufferDesc.nextInChain = nullptr;
bufferDesc.label = label;
bufferDesc.usage = usage;
bufferDesc.size = size;
bufferDesc.mappedAtCreation = false;
return wgpuDeviceCreateBuffer(device, &bufferDesc);
const WGPUTextureDescriptor textureDesc {
.usage = WGPUTextureUsage_RenderAttachment,
.dimension = WGPUTextureDimension_2D, .size = { width, height, 1 },
.format = format, .mipLevelCount = 1, .sampleCount = sc
};
return wgpuDeviceCreateTexture(device, &textureDesc);
}
void WgContext::releaseSampler(WGPUSampler& sampler)
WGPUTextureView WgContext::createTextureView(WGPUTexture texture)
{
if (sampler) {
wgpuSamplerRelease(sampler);
sampler = nullptr;
const WGPUTextureViewDescriptor textureViewDesc {
.format = wgpuTextureGetFormat(texture),
.dimension = WGPUTextureViewDimension_2D,
.baseMipLevel = 0,
.mipLevelCount = 1,
.baseArrayLayer = 0,
.arrayLayerCount = 1,
.aspect = WGPUTextureAspect_All
};
return wgpuTextureCreateView(texture, &textureViewDesc);
}
void WgContext::releaseTextureView(WGPUTextureView& textureView)
{
if (textureView) {
wgpuTextureViewRelease(textureView);
textureView = nullptr;
}
}
@ -246,15 +164,89 @@ void WgContext::releaseTexture(WGPUTexture& texture)
}
void WgContext::releaseTextureView(WGPUTextureView& textureView)
void WgContext::releaseSampler(WGPUSampler& sampler)
{
if (textureView) {
wgpuTextureViewRelease(textureView);
textureView = nullptr;
if (sampler) {
wgpuSamplerRelease(sampler);
sampler = nullptr;
}
}
bool WgContext::allocateBufferUniform(WGPUBuffer& buffer, const void* data, uint64_t size)
{
if ((buffer) && (wgpuBufferGetSize(buffer) >= size))
wgpuQueueWriteBuffer(queue, buffer, 0, data, size);
else {
releaseBuffer(buffer);
const WGPUBufferDescriptor bufferDesc { .usage = WGPUBufferUsage_CopyDst | WGPUBufferUsage_Uniform, .size = size };
buffer = wgpuDeviceCreateBuffer(device, &bufferDesc);
wgpuQueueWriteBuffer(queue, buffer, 0, data, size);
return true;
}
return false;
}
bool WgContext::allocateBufferVertex(WGPUBuffer& buffer, const float* data, uint64_t size)
{
if ((buffer) && (wgpuBufferGetSize(buffer) >= size))
wgpuQueueWriteBuffer(queue, buffer, 0, data, size);
else {
releaseBuffer(buffer);
const WGPUBufferDescriptor bufferDesc {
.usage = WGPUBufferUsage_CopyDst | WGPUBufferUsage_Vertex,
.size = size > WG_VERTEX_BUFFER_MIN_SIZE ? size : WG_VERTEX_BUFFER_MIN_SIZE
};
buffer = wgpuDeviceCreateBuffer(device, &bufferDesc);
wgpuQueueWriteBuffer(queue, buffer, 0, data, size);
return true;
}
return false;
}
bool WgContext::allocateBufferIndex(WGPUBuffer& buffer, const uint32_t* data, uint64_t size)
{
if ((buffer) && (wgpuBufferGetSize(buffer) >= size))
wgpuQueueWriteBuffer(queue, buffer, 0, data, size);
else {
releaseBuffer(buffer);
const WGPUBufferDescriptor bufferDesc {
.usage = WGPUBufferUsage_CopyDst | WGPUBufferUsage_Index,
.size = size > WG_INDEX_BUFFER_MIN_SIZE ? size : WG_INDEX_BUFFER_MIN_SIZE
};
buffer = wgpuDeviceCreateBuffer(device, &bufferDesc);
wgpuQueueWriteBuffer(queue, buffer, 0, data, size);
return true;
}
return false;
}
bool WgContext::allocateBufferIndexFan(uint64_t vertexCount)
{
uint64_t indexCount = (vertexCount - 2) * 3;
if ((!bufferIndexFan) || (wgpuBufferGetSize(bufferIndexFan) < indexCount * sizeof(uint32_t))) {
tvg::Array<uint32_t> indexes(indexCount);
for (size_t i = 0; i < vertexCount - 2; i++) {
indexes.push(0);
indexes.push(i + 1);
indexes.push(i + 2);
}
releaseBuffer(bufferIndexFan);
WGPUBufferDescriptor bufferDesc{
.usage = WGPUBufferUsage_CopyDst | WGPUBufferUsage_Index,
.size = indexCount * sizeof(uint32_t)
};
bufferIndexFan = wgpuDeviceCreateBuffer(device, &bufferDesc);
wgpuQueueWriteBuffer(queue, bufferIndexFan, 0, &indexes[0], indexCount * sizeof(uint32_t));
return true;
}
return false;
}
void WgContext::releaseBuffer(WGPUBuffer& buffer)
{
if (buffer) {
@ -263,540 +255,3 @@ void WgContext::releaseBuffer(WGPUBuffer& buffer)
buffer = nullptr;
}
}
void WgContext::allocateVertexBuffer(WGPUBuffer& buffer, const void *data, uint64_t size)
{
if ((buffer) && (wgpuBufferGetSize(buffer) >= size))
wgpuQueueWriteBuffer(queue, buffer, 0, data, size);
else {
// create new buffer and upload data
releaseBuffer(buffer);
WGPUBufferDescriptor bufferDesc{};
bufferDesc.nextInChain = nullptr;
bufferDesc.label = "The vertex buffer";
bufferDesc.usage = WGPUBufferUsage_CopyDst | WGPUBufferUsage_Vertex;
bufferDesc.size = size > WG_VERTEX_BUFFER_MIN_SIZE ? size : WG_VERTEX_BUFFER_MIN_SIZE;
bufferDesc.mappedAtCreation = false;
buffer = wgpuDeviceCreateBuffer(device, &bufferDesc);
wgpuQueueWriteBuffer(queue, buffer, 0, data, size);
}
}
void WgContext::allocateIndexBuffer(WGPUBuffer& buffer, const void *data, uint64_t size)
{
if ((buffer) && (wgpuBufferGetSize(buffer) >= size))
wgpuQueueWriteBuffer(queue, buffer, 0, data, size);
else {
// create new buffer and upload data
releaseBuffer(buffer);
WGPUBufferDescriptor bufferDesc{};
bufferDesc.nextInChain = nullptr;
bufferDesc.label = "The index buffer";
bufferDesc.usage = WGPUBufferUsage_CopyDst | WGPUBufferUsage_Index;
bufferDesc.size = size > WG_INDEX_BUFFER_MIN_SIZE ? size : WG_INDEX_BUFFER_MIN_SIZE;
bufferDesc.mappedAtCreation = false;
buffer = wgpuDeviceCreateBuffer(device, &bufferDesc);
wgpuQueueWriteBuffer(queue, buffer, 0, data, size);
}
}
void WgContext::allocateIndexBufferFan(uint64_t vertsCount)
{
assert(vertsCount >= 3);
uint64_t indexCount = (vertsCount - 2) * 3;
if ((!indexBufferFan) || (wgpuBufferGetSize(indexBufferFan) < indexCount * sizeof(uint32_t))) {
Array<uint32_t> indexes(indexCount);
for (size_t i = 0; i < vertsCount - 2; i++) {
indexes.push(0);
indexes.push(i + 1);
indexes.push(i + 2);
}
releaseBuffer(indexBufferFan);
WGPUBufferDescriptor bufferDesc{};
bufferDesc.nextInChain = nullptr;
bufferDesc.label = "The index buffer";
bufferDesc.usage = WGPUBufferUsage_CopyDst | WGPUBufferUsage_Index;
bufferDesc.size = indexCount * sizeof(uint32_t);
bufferDesc.mappedAtCreation = false;
indexBufferFan = wgpuDeviceCreateBuffer(device, &bufferDesc);
wgpuQueueWriteBuffer(queue, indexBufferFan, 0, &indexes[0], indexCount * sizeof(uint32_t));
}
}
void WgContext::releaseVertexBuffer(WGPUBuffer& buffer)
{
releaseBuffer(buffer);
}
void WgContext::releaseIndexBuffer(WGPUBuffer& buffer)
{
releaseBuffer(buffer);
}
//*****************************************************************************
// bind group
//*****************************************************************************
void WgBindGroup::set(WGPURenderPassEncoder encoder, uint32_t groupIndex)
{
wgpuRenderPassEncoderSetBindGroup(encoder, groupIndex, mBindGroup, 0, nullptr);
}
void WgBindGroup::set(WGPUComputePassEncoder encoder, uint32_t groupIndex)
{
wgpuComputePassEncoderSetBindGroup(encoder, groupIndex, mBindGroup, 0, nullptr);
}
WGPUBindGroupEntry WgBindGroup::makeBindGroupEntryBuffer(uint32_t binding, WGPUBuffer buffer)
{
WGPUBindGroupEntry bindGroupEntry{};
bindGroupEntry.nextInChain = nullptr;
bindGroupEntry.binding = binding;
bindGroupEntry.buffer = buffer;
bindGroupEntry.offset = 0;
bindGroupEntry.size = wgpuBufferGetSize(buffer);
bindGroupEntry.sampler = nullptr;
bindGroupEntry.textureView = nullptr;
return bindGroupEntry;
}
WGPUBindGroupEntry WgBindGroup::makeBindGroupEntrySampler(uint32_t binding, WGPUSampler sampler)
{
WGPUBindGroupEntry bindGroupEntry{};
bindGroupEntry.nextInChain = nullptr;
bindGroupEntry.binding = binding;
bindGroupEntry.buffer = nullptr;
bindGroupEntry.offset = 0;
bindGroupEntry.size = 0;
bindGroupEntry.sampler = sampler;
bindGroupEntry.textureView = nullptr;
return bindGroupEntry;
}
WGPUBindGroupEntry WgBindGroup::makeBindGroupEntryTextureView(uint32_t binding, WGPUTextureView textureView)
{
WGPUBindGroupEntry bindGroupEntry{};
bindGroupEntry.nextInChain = nullptr;
bindGroupEntry.binding = binding;
bindGroupEntry.buffer = nullptr;
bindGroupEntry.offset = 0;
bindGroupEntry.size = 0;
bindGroupEntry.sampler = nullptr;
bindGroupEntry.textureView = textureView;
return bindGroupEntry;
}
WGPUBindGroupLayoutEntry WgBindGroup::makeBindGroupLayoutEntryBuffer(uint32_t binding)
{
WGPUBindGroupLayoutEntry bindGroupLayoutEntry{};
bindGroupLayoutEntry.nextInChain = nullptr;
bindGroupLayoutEntry.binding = binding;
bindGroupLayoutEntry.visibility = WGPUShaderStage_Vertex | WGPUShaderStage_Fragment | WGPUShaderStage_Compute;
bindGroupLayoutEntry.buffer.nextInChain = nullptr;
bindGroupLayoutEntry.buffer.type = WGPUBufferBindingType_Uniform;
bindGroupLayoutEntry.buffer.hasDynamicOffset = false;
bindGroupLayoutEntry.buffer.minBindingSize = 0;
return bindGroupLayoutEntry;
}
WGPUBindGroupLayoutEntry WgBindGroup::makeBindGroupLayoutEntrySampler(uint32_t binding)
{
WGPUBindGroupLayoutEntry bindGroupLayoutEntry{};
bindGroupLayoutEntry.nextInChain = nullptr;
bindGroupLayoutEntry.binding = binding;
bindGroupLayoutEntry.visibility = WGPUShaderStage_Fragment;
bindGroupLayoutEntry.sampler.nextInChain = nullptr;
bindGroupLayoutEntry.sampler.type = WGPUSamplerBindingType_Filtering;
return bindGroupLayoutEntry;
}
WGPUBindGroupLayoutEntry WgBindGroup::makeBindGroupLayoutEntryTexture(uint32_t binding)
{
WGPUBindGroupLayoutEntry bindGroupLayoutEntry{};
bindGroupLayoutEntry.nextInChain = nullptr;
bindGroupLayoutEntry.binding = binding;
bindGroupLayoutEntry.visibility = WGPUShaderStage_Fragment | WGPUShaderStage_Compute;
bindGroupLayoutEntry.texture.nextInChain = nullptr;
bindGroupLayoutEntry.texture.sampleType = WGPUTextureSampleType_Float;
bindGroupLayoutEntry.texture.viewDimension = WGPUTextureViewDimension_2D;
bindGroupLayoutEntry.texture.multisampled = false;
return bindGroupLayoutEntry;
}
WGPUBindGroupLayoutEntry WgBindGroup::makeBindGroupLayoutEntryStorage(uint32_t binding, WGPUStorageTextureAccess access, WGPUTextureFormat format)
{
WGPUBindGroupLayoutEntry bindGroupLayoutEntry{};
bindGroupLayoutEntry.nextInChain = nullptr;
bindGroupLayoutEntry.binding = binding;
bindGroupLayoutEntry.visibility = WGPUShaderStage_Fragment | WGPUShaderStage_Compute;
bindGroupLayoutEntry.storageTexture.nextInChain = nullptr;
bindGroupLayoutEntry.storageTexture.access = access;
bindGroupLayoutEntry.storageTexture.format = format;
bindGroupLayoutEntry.storageTexture.viewDimension = WGPUTextureViewDimension_2D;
return bindGroupLayoutEntry;
}
WGPUBuffer WgBindGroup::createBuffer(WGPUDevice device, WGPUQueue queue, const void *data, size_t size)
{
WGPUBufferDescriptor bufferDescriptor{};
bufferDescriptor.nextInChain = nullptr;
bufferDescriptor.label = "The uniform buffer";
bufferDescriptor.usage = WGPUBufferUsage_CopyDst | WGPUBufferUsage_Uniform;
bufferDescriptor.size = size;
bufferDescriptor.mappedAtCreation = false;
WGPUBuffer buffer = wgpuDeviceCreateBuffer(device, &bufferDescriptor);
assert(buffer);
wgpuQueueWriteBuffer(queue, buffer, 0, data, size);
return buffer;
}
WGPUBindGroup WgBindGroup::createBindGroup(WGPUDevice device, WGPUBindGroupLayout layout, const WGPUBindGroupEntry* bindGroupEntries, uint32_t count)
{
WGPUBindGroupDescriptor bindGroupDesc{};
bindGroupDesc.nextInChain = nullptr;
bindGroupDesc.label = "The binding group";
bindGroupDesc.layout = layout;
bindGroupDesc.entryCount = count;
bindGroupDesc.entries = bindGroupEntries;
return wgpuDeviceCreateBindGroup(device, &bindGroupDesc);
}
WGPUBindGroupLayout WgBindGroup::createBindGroupLayout(WGPUDevice device, const WGPUBindGroupLayoutEntry* bindGroupLayoutEntries, uint32_t count)
{
WGPUBindGroupLayoutDescriptor bindGroupLayoutDesc{};
bindGroupLayoutDesc.nextInChain = nullptr;
bindGroupLayoutDesc.label = "The bind group layout";
bindGroupLayoutDesc.entryCount = count;
bindGroupLayoutDesc.entries = bindGroupLayoutEntries; // @binding
return wgpuDeviceCreateBindGroupLayout(device, &bindGroupLayoutDesc);
}
void WgBindGroup::releaseBuffer(WGPUBuffer& buffer)
{
if (buffer) {
wgpuBufferDestroy(buffer);
wgpuBufferRelease(buffer);
buffer = nullptr;
}
}
void WgBindGroup::releaseBindGroup(WGPUBindGroup& bindGroup)
{
if (bindGroup) wgpuBindGroupRelease(bindGroup);
bindGroup = nullptr;
}
void WgBindGroup::releaseBindGroupLayout(WGPUBindGroupLayout& bindGroupLayout)
{
if (bindGroupLayout) wgpuBindGroupLayoutRelease(bindGroupLayout);
bindGroupLayout = nullptr;
}
//*****************************************************************************
// pipeline
//*****************************************************************************
void WgPipeline::release()
{
destroyShaderModule(mShaderModule);
destroyPipelineLayout(mPipelineLayout);
}
WGPUPipelineLayout WgPipeline::createPipelineLayout(WGPUDevice device, const WGPUBindGroupLayout* bindGroupLayouts, uint32_t count)
{
WGPUPipelineLayoutDescriptor pipelineLayoutDesc{};
pipelineLayoutDesc.nextInChain = nullptr;
pipelineLayoutDesc.label = "The pipeline layout";
pipelineLayoutDesc.bindGroupLayoutCount = count;
pipelineLayoutDesc.bindGroupLayouts = bindGroupLayouts;
return wgpuDeviceCreatePipelineLayout(device, &pipelineLayoutDesc);
}
WGPUShaderModule WgPipeline::createShaderModule(WGPUDevice device, const char* code, const char* label)
{
WGPUShaderModuleWGSLDescriptor shaderModuleWGSLDesc{};
shaderModuleWGSLDesc.chain.next = nullptr;
shaderModuleWGSLDesc.chain.sType = WGPUSType_ShaderModuleWGSLDescriptor;
shaderModuleWGSLDesc.code = code;
WGPUShaderModuleDescriptor shaderModuleDesc{};
shaderModuleDesc.nextInChain = &shaderModuleWGSLDesc.chain;
shaderModuleDesc.label = label;
#ifndef __EMSCRIPTEN__
shaderModuleDesc.hintCount = 0;
shaderModuleDesc.hints = nullptr;
#endif
return wgpuDeviceCreateShaderModule(device, &shaderModuleDesc);
}
void WgPipeline::destroyPipelineLayout(WGPUPipelineLayout& pipelineLayout)
{
if (pipelineLayout) wgpuPipelineLayoutRelease(pipelineLayout);
pipelineLayout = nullptr;
}
void WgPipeline::destroyShaderModule(WGPUShaderModule& shaderModule)
{
if (shaderModule) wgpuShaderModuleRelease(shaderModule);
shaderModule = nullptr;
}
//*****************************************************************************
// render pipeline
//*****************************************************************************
void WgRenderPipeline::allocate(WGPUDevice device, WgPipelineBlendType blendType, WGPUColorWriteMask writeMask,
WGPUVertexBufferLayout vertexBufferLayouts[], uint32_t attribsCount,
WGPUBindGroupLayout bindGroupLayouts[], uint32_t bindGroupsCount,
WGPUCompareFunction compareFront, WGPUStencilOperation operationFront,
WGPUCompareFunction compareBack, WGPUStencilOperation operationBack,
const char* shaderSource, const char* shaderLabel, const char* pipelineLabel)
{
mShaderModule = createShaderModule(device, shaderSource, shaderLabel);
assert(mShaderModule);
mPipelineLayout = createPipelineLayout(device, bindGroupLayouts, bindGroupsCount);
assert(mPipelineLayout);
mRenderPipeline = createRenderPipeline(device, blendType, writeMask,
vertexBufferLayouts, attribsCount,
compareFront, operationFront,
compareBack, operationBack,
mPipelineLayout, mShaderModule, pipelineLabel);
assert(mRenderPipeline);
}
void WgRenderPipeline::release()
{
destroyRenderPipeline(mRenderPipeline);
WgPipeline::release();
}
void WgRenderPipeline::set(WGPURenderPassEncoder renderPassEncoder)
{
wgpuRenderPassEncoderSetPipeline(renderPassEncoder, mRenderPipeline);
};
WGPUBlendState WgRenderPipeline::makeBlendState(WgPipelineBlendType blendType)
{
WGPUBlendState blendState{};
// srcover
if (blendType == WgPipelineBlendType::SrcOver) {
blendState.color.operation = WGPUBlendOperation_Add;
blendState.color.srcFactor = WGPUBlendFactor_One;
blendState.color.dstFactor = WGPUBlendFactor_Zero;
// normal
} else if (blendType == WgPipelineBlendType::Normal) {
blendState.color.operation = WGPUBlendOperation_Add;
blendState.color.srcFactor = WGPUBlendFactor_One;
blendState.color.dstFactor = WGPUBlendFactor_OneMinusSrcAlpha;
// custom
} else if (blendType == WgPipelineBlendType::Custom) {
blendState.color.operation = WGPUBlendOperation_Add;
blendState.color.srcFactor = WGPUBlendFactor_One;
blendState.color.dstFactor = WGPUBlendFactor_Zero;
}
blendState.alpha = blendState.color;
return blendState;
}
WGPUColorTargetState WgRenderPipeline::makeColorTargetState(const WGPUBlendState* blendState, const WGPUColorWriteMask writeMask)
{
WGPUColorTargetState colorTargetState{};
colorTargetState.nextInChain = nullptr;
colorTargetState.format = WGPUTextureFormat_RGBA8Unorm;
colorTargetState.blend = blendState;
colorTargetState.writeMask = writeMask;
return colorTargetState;
}
WGPUVertexBufferLayout WgRenderPipeline::makeVertexBufferLayout(const WGPUVertexAttribute* vertexAttributes, uint32_t count, uint64_t stride)
{
WGPUVertexBufferLayout vertexBufferLayoutPos{};
vertexBufferLayoutPos.arrayStride = stride;
vertexBufferLayoutPos.stepMode = WGPUVertexStepMode_Vertex;
vertexBufferLayoutPos.attributeCount = count;
vertexBufferLayoutPos.attributes = vertexAttributes;
return vertexBufferLayoutPos;
}
WGPUVertexState WgRenderPipeline::makeVertexState(WGPUShaderModule shaderModule, const WGPUVertexBufferLayout* buffers, uint32_t count)
{
WGPUVertexState vertexState{};
vertexState.nextInChain = nullptr;
vertexState.module = shaderModule;
vertexState.entryPoint = "vs_main";
vertexState.constantCount = 0;
vertexState.constants = nullptr;
vertexState.bufferCount = count;
vertexState.buffers = buffers;
return vertexState;
}
WGPUPrimitiveState WgRenderPipeline::makePrimitiveState()
{
WGPUPrimitiveState primitiveState{};
primitiveState.nextInChain = nullptr;
primitiveState.topology = WGPUPrimitiveTopology_TriangleList;
primitiveState.stripIndexFormat = WGPUIndexFormat_Undefined;
primitiveState.frontFace = WGPUFrontFace_CCW;
primitiveState.cullMode = WGPUCullMode_None;
return primitiveState;
}
WGPUDepthStencilState WgRenderPipeline::makeDepthStencilState(WGPUCompareFunction compareFront, WGPUStencilOperation operationFront, WGPUCompareFunction compareBack, WGPUStencilOperation operationBack)
{
WGPUDepthStencilState depthStencilState{};
depthStencilState.nextInChain = nullptr;
depthStencilState.format = WGPUTextureFormat_Stencil8;
depthStencilState.depthWriteEnabled = false;
depthStencilState.depthCompare = WGPUCompareFunction_Always;
depthStencilState.stencilFront.compare = compareFront;
depthStencilState.stencilFront.failOp = operationFront;
depthStencilState.stencilFront.depthFailOp = operationFront;
depthStencilState.stencilFront.passOp = operationFront;
depthStencilState.stencilBack.compare = compareBack;
depthStencilState.stencilBack.failOp = operationBack;
depthStencilState.stencilBack.depthFailOp = operationBack;
depthStencilState.stencilBack.passOp = operationBack;
depthStencilState.stencilReadMask = 0xFFFFFFFF;
depthStencilState.stencilWriteMask = 0xFFFFFFFF;
depthStencilState.depthBias = 0;
depthStencilState.depthBiasSlopeScale = 0.0f;
depthStencilState.depthBiasClamp = 0.0f;
return depthStencilState;
}
WGPUMultisampleState WgRenderPipeline::makeMultisampleState()
{
WGPUMultisampleState multisampleState{};
multisampleState.nextInChain = nullptr;
multisampleState.count = 1;
multisampleState.mask = 0xFFFFFFFF;
multisampleState.alphaToCoverageEnabled = false;
return multisampleState;
}
WGPUFragmentState WgRenderPipeline::makeFragmentState(WGPUShaderModule shaderModule, WGPUColorTargetState* targets, uint32_t size)
{
WGPUFragmentState fragmentState{};
fragmentState.nextInChain = nullptr;
fragmentState.module = shaderModule;
fragmentState.entryPoint = "fs_main";
fragmentState.constantCount = 0;
fragmentState.constants = nullptr;
fragmentState.targetCount = size;
fragmentState.targets = targets;
return fragmentState;
}
WGPURenderPipeline WgRenderPipeline::createRenderPipeline(WGPUDevice device, WgPipelineBlendType blendType, WGPUColorWriteMask writeMask,
WGPUVertexBufferLayout vertexBufferLayouts[], uint32_t attribsCount,
WGPUCompareFunction compareFront, WGPUStencilOperation operationFront,
WGPUCompareFunction compareBack, WGPUStencilOperation operationBack,
WGPUPipelineLayout pipelineLayout, WGPUShaderModule shaderModule,
const char* pipelineName)
{
WGPUBlendState blendState = makeBlendState(blendType);
WGPUColorTargetState colorTargetStates[] = {
makeColorTargetState(&blendState, writeMask)
};
WGPUVertexState vertexState = makeVertexState(shaderModule, vertexBufferLayouts, attribsCount);
WGPUPrimitiveState primitiveState = makePrimitiveState();
WGPUDepthStencilState depthStencilState = makeDepthStencilState(compareFront, operationFront, compareBack, operationBack);
WGPUMultisampleState multisampleState = makeMultisampleState();
WGPUFragmentState fragmentState = makeFragmentState(shaderModule, colorTargetStates, 1);
WGPURenderPipelineDescriptor renderPipelineDesc{};
renderPipelineDesc.nextInChain = nullptr;
renderPipelineDesc.label = pipelineName;
renderPipelineDesc.layout = pipelineLayout;
renderPipelineDesc.vertex = vertexState;
renderPipelineDesc.primitive = primitiveState;
renderPipelineDesc.depthStencil = &depthStencilState;
renderPipelineDesc.multisample = multisampleState;
renderPipelineDesc.fragment = &fragmentState;
return wgpuDeviceCreateRenderPipeline(device, &renderPipelineDesc);
}
void WgRenderPipeline::destroyRenderPipeline(WGPURenderPipeline& renderPipeline)
{
if (renderPipeline) wgpuRenderPipelineRelease(renderPipeline);
renderPipeline = nullptr;
}
//*****************************************************************************
// compute pipeline
//*****************************************************************************
void WgComputePipeline::allocate(WGPUDevice device,
WGPUBindGroupLayout bindGroupLayouts[], uint32_t bindGroupsCount,
const char* shaderSource, const char* shaderLabel, const char* pipelineLabel)
{
mShaderModule = createShaderModule(device, shaderSource, shaderLabel);
assert(mShaderModule);
mPipelineLayout = createPipelineLayout(device, bindGroupLayouts, bindGroupsCount);
assert(mPipelineLayout);
WGPUComputePipelineDescriptor computePipelineDesc{};
computePipelineDesc.nextInChain = nullptr;
computePipelineDesc.label = pipelineLabel;
computePipelineDesc.layout = mPipelineLayout;
computePipelineDesc.compute.nextInChain = nullptr;
computePipelineDesc.compute.module = mShaderModule;
computePipelineDesc.compute.entryPoint = "cs_main";
computePipelineDesc.compute.constantCount = 0;
computePipelineDesc.compute.constants = nullptr;
mComputePipeline = wgpuDeviceCreateComputePipeline(device, &computePipelineDesc);
assert(mComputePipeline);
}
void WgComputePipeline::release()
{
if (mComputePipeline) wgpuComputePipelineRelease(mComputePipeline);
mComputePipeline = nullptr;
WgPipeline::release();
}
void WgComputePipeline::set(WGPUComputePassEncoder computePassEncoder)
{
wgpuComputePassEncoderSetPipeline(computePassEncoder, mComputePipeline);
}

150
src/renderer/wg_engine/tvgWgCommon.h Normal file → Executable file
View file

@ -25,151 +25,51 @@
#include <cassert>
#include <webgpu/webgpu.h>
#include "tvgCommon.h"
#include "tvgRender.h"
#define WG_VERTEX_BUFFER_MIN_SIZE 2048
#define WG_INDEX_BUFFER_MIN_SIZE 2048
struct WgCompositor: public Compositor {
BlendMethod blendMethod;
};
enum class WgPipelineBlendType {
SrcOver = 0,
Normal,
Custom
};
struct WgPipelines;
class WgPipelines;
struct WgContext {
// external webgpu handles
WGPUInstance instance{};
WGPUSurface surface{};
// common webgpu handles
WGPUAdapter adapter{};
WGPUDevice device{};
WGPUQueue queue{};
WGPUFeatureName featureNames[32]{};
WGPUAdapterProperties adapterProperties{};
WGPUSupportedLimits supportedLimits{};
WGPUTextureFormat preferredFormat{};
// external handles (do not release)
WgPipelines* pipelines{};
// shared webgpu assets
WGPUBuffer bufferIndexFan{};
WGPUSampler samplerNearest{};
WGPUSampler samplerLinear{};
WGPUBuffer indexBufferFan{};
WgPipelines* pipelines{}; // external handle (do not release)
void initialize(WGPUInstance instance, WGPUSurface surface);
void release();
void executeCommandEncoder(WGPUCommandEncoder commandEncoder);
// create common objects
WGPUSampler createSampler(WGPUFilterMode filter, WGPUMipmapFilterMode mipmapFilter);
WGPUTexture createTexture(uint32_t width, uint32_t height, WGPUTextureFormat format);
WGPUTexture createTexStorage(uint32_t width, uint32_t height, WGPUTextureFormat format, uint32_t sc = 1);
WGPUTexture createTexStencil(uint32_t width, uint32_t height, WGPUTextureFormat format, uint32_t sc = 1);
WGPUTextureView createTextureView(WGPUTexture texture);
WGPUSampler createSampler(WGPUFilterMode minFilter, WGPUMipmapFilterMode mipmapFilter);
WGPUTexture createTexture2d(WGPUTextureUsageFlags usage, WGPUTextureFormat format, uint32_t width, uint32_t height, char const * label);
WGPUTexture createTexture2dMS(WGPUTextureUsageFlags usage, WGPUTextureFormat format, uint32_t width, uint32_t height, uint32_t sc, char const * label);
WGPUTextureView createTextureView2d(WGPUTexture texture, char const * label);
WGPUBuffer createBuffer(WGPUBufferUsageFlags usage, uint64_t size, char const * label);
void releaseSampler(WGPUSampler& sampler);
void releaseTexture(WGPUTexture& texture);
// release common objects
void releaseTextureView(WGPUTextureView& textureView);
void releaseTexture(WGPUTexture& texture);
void releaseSampler(WGPUSampler& sampler);
// create buffer objects (return true, if buffer handle was changed)
bool allocateBufferUniform(WGPUBuffer& buffer, const void* data, uint64_t size);
bool allocateBufferVertex(WGPUBuffer& buffer, const float* data, uint64_t size);
bool allocateBufferIndex(WGPUBuffer& buffer, const uint32_t* data, uint64_t size);
bool allocateBufferIndexFan(uint64_t vertexCount);
// release buffer objects
void releaseBuffer(WGPUBuffer& buffer);
void allocateVertexBuffer(WGPUBuffer& buffer, const void *data, uint64_t size);
void allocateIndexBuffer(WGPUBuffer& buffer, const void *data, uint64_t size);
void allocateIndexBufferFan(uint64_t size);
void releaseVertexBuffer(WGPUBuffer& buffer);
void releaseIndexBuffer(WGPUBuffer& buffer);
};
struct WgBindGroup
{
WGPUBindGroup mBindGroup{};
void set(WGPURenderPassEncoder encoder, uint32_t groupIndex);
void set(WGPUComputePassEncoder encoder, uint32_t groupIndex);
static WGPUBindGroupEntry makeBindGroupEntryBuffer(uint32_t binding, WGPUBuffer buffer);
static WGPUBindGroupEntry makeBindGroupEntrySampler(uint32_t binding, WGPUSampler sampler);
static WGPUBindGroupEntry makeBindGroupEntryTextureView(uint32_t binding, WGPUTextureView textureView);
static WGPUBindGroupLayoutEntry makeBindGroupLayoutEntryBuffer(uint32_t binding);
static WGPUBindGroupLayoutEntry makeBindGroupLayoutEntrySampler(uint32_t binding);
static WGPUBindGroupLayoutEntry makeBindGroupLayoutEntryTexture(uint32_t binding);
static WGPUBindGroupLayoutEntry makeBindGroupLayoutEntryStorage(uint32_t binding, WGPUStorageTextureAccess access, WGPUTextureFormat format);
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 WGPUBindGroupLayout createBindGroupLayout(WGPUDevice device, const WGPUBindGroupLayoutEntry* bindGroupLayoutEntries, uint32_t count);
static void releaseBuffer(WGPUBuffer& buffer);
static void releaseBindGroup(WGPUBindGroup& bindGroup);
static void releaseBindGroupLayout(WGPUBindGroupLayout& bindGroupLayout);
};
struct WgPipeline
{
protected:
WGPUPipelineLayout mPipelineLayout{};
WGPUShaderModule mShaderModule{};
public:
virtual ~WgPipeline() {}
virtual void initialize(WGPUDevice device) = 0;
virtual void release();
static WGPUPipelineLayout createPipelineLayout(WGPUDevice device, const WGPUBindGroupLayout* bindGroupLayouts, uint32_t count);
static WGPUShaderModule createShaderModule(WGPUDevice device, const char* code, const char* label);
static void destroyPipelineLayout(WGPUPipelineLayout& pipelineLayout);
static void destroyShaderModule(WGPUShaderModule& shaderModule);
};
struct WgRenderPipeline: public WgPipeline
{
protected:
WGPURenderPipeline mRenderPipeline{};
void allocate(WGPUDevice device, WgPipelineBlendType blendType, WGPUColorWriteMask writeMask,
WGPUVertexBufferLayout vertexBufferLayouts[], uint32_t attribsCount,
WGPUBindGroupLayout bindGroupLayouts[], uint32_t bindGroupsCount,
WGPUCompareFunction compareFront, WGPUStencilOperation operationFront,
WGPUCompareFunction compareBack, WGPUStencilOperation operationBack,
const char* shaderSource, const char* shaderLabel, const char* pipelineLabel);
public:
void release() override;
void set(WGPURenderPassEncoder renderPassEncoder);
static WGPUBlendState makeBlendState(WgPipelineBlendType blendType);
static WGPUColorTargetState makeColorTargetState(const WGPUBlendState* blendState, const WGPUColorWriteMask writeMask);
static WGPUVertexBufferLayout makeVertexBufferLayout(const WGPUVertexAttribute* vertexAttributes, uint32_t count, uint64_t stride);
static WGPUVertexState makeVertexState(WGPUShaderModule shaderModule, const WGPUVertexBufferLayout* buffers, uint32_t count);
static WGPUPrimitiveState makePrimitiveState();
static WGPUDepthStencilState makeDepthStencilState(WGPUCompareFunction compareFront, WGPUStencilOperation operationFront, WGPUCompareFunction compareBack, WGPUStencilOperation operationBack);
static WGPUMultisampleState makeMultisampleState();
static WGPUFragmentState makeFragmentState(WGPUShaderModule shaderModule, WGPUColorTargetState* targets, uint32_t size);
static WGPURenderPipeline createRenderPipeline(WGPUDevice device, WgPipelineBlendType blendType, WGPUColorWriteMask writeMask,
WGPUVertexBufferLayout vertexBufferLayouts[], uint32_t attribsCount,
WGPUCompareFunction compareFront, WGPUStencilOperation operationFront,
WGPUCompareFunction compareBack, WGPUStencilOperation operationBack,
WGPUPipelineLayout pipelineLayout, WGPUShaderModule shaderModule,
const char* pipelineLabel);
static void destroyRenderPipeline(WGPURenderPipeline& renderPipeline);
};
#define WG_COMPUTE_WORKGROUP_SIZE_X 8
#define WG_COMPUTE_WORKGROUP_SIZE_Y 8
struct WgComputePipeline: public WgPipeline
{
protected:
WGPUComputePipeline mComputePipeline{};
void allocate(WGPUDevice device,
WGPUBindGroupLayout bindGroupLayouts[], uint32_t bindGroupsCount,
const char* shaderSource, const char* shaderLabel, const char* pipelineLabel);
public:
void release() override;
void set(WGPUComputePassEncoder computePassEncoder);
};
#endif // _TVG_WG_COMMON_H_

509
src/renderer/wg_engine/tvgWgCompositor.cpp Executable file
View file

@ -0,0 +1,509 @@
/*
* Copyright (c) 2024 the ThorVG project. All rights reserved.
* Permission is hereby granted, free of charge, to any person obtaining a copy
* of this software and associated documentation files (the "Software"), to deal
* in the Software without restriction, including without limitation the rights
* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
* copies of the Software, and to permit persons to whom the Software is
* furnished to do so, subject to the following conditions:
* The above copyright notice and this permission notice shall be included in all
* copies or substantial portions of the Software.
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
* SOFTWARE.
*/
#include "tvgWgCompositor.h"
#include "tvgWgShaderTypes.h"
void WgCompositor::initialize(WgContext& context, uint32_t width, uint32_t height)
{
// pipelines (external handle, do not release)
pipelines = context.pipelines;
// store render target dimensions
this->width = width;
this->height = height;
// allocate global stencil buffer handles
texStencil = context.createTexStencil(width, height, WGPUTextureFormat_Stencil8);
texViewStencil = context.createTextureView(texStencil);
// allocate global view matrix handles
WgShaderTypeMat4x4f viewMat(width, height);
context.allocateBufferUniform(bufferViewMat, &viewMat, sizeof(viewMat));
bindGroupViewMat = pipelines->layouts.createBindGroupBuffer1Un(bufferViewMat);
// initialize opacity pool
for (uint32_t i = 0; i < 256; i++) {
float opacity = i / 255.0f;
context.allocateBufferUniform(bufferOpacities[i], &opacity, sizeof(float));
bindGroupOpacities[i] = pipelines->layouts.createBindGroupBuffer1Un(bufferOpacities[i]);
}
// initialize intermediate render storages
storageClipPath.initialize(context, width, height);
storageInterm.initialize(context, width, height);
storageDstCopy.initialize(context, width, height);
}
void WgCompositor::release(WgContext& context)
{
// release intermediate render storages
storageInterm.release(context);
storageDstCopy.release(context);
storageClipPath.release(context);
// release opacity pool
for (uint32_t i = 0; i < 256; i++) {
context.pipelines->layouts.releaseBindGroup(bindGroupOpacities[i]);
context.releaseBuffer(bufferOpacities[i]);
}
// release global view matrix handles
context.pipelines->layouts.releaseBindGroup(bindGroupViewMat);
context.releaseBuffer(bufferViewMat);
// release global stencil buffer handles
context.releaseTextureView(texViewStencil);
context.releaseTexture(texStencil);
height = 0;
width = 0;
pipelines = nullptr;
}
WgPipelineBlendType WgCompositor::blendMethodToBlendType(BlendMethod blendMethod)
{
if (blendMethod == BlendMethod::Normal) return WgPipelineBlendType::Normal;
if (blendMethod == BlendMethod::SrcOver) return WgPipelineBlendType::SrcOver;
return WgPipelineBlendType::Custom;
}
void WgCompositor::beginRenderPass(WGPUCommandEncoder commandEncoder, WgRenderStorage* target, bool clear)
{
assert(commandEncoder);
assert(target);
this->currentTarget = target;
this->commandEncoder = commandEncoder;
WGPURenderPassDepthStencilAttachment depthStencilAttachment{ .view = texViewStencil, .stencilLoadOp = WGPULoadOp_Clear, .stencilStoreOp = WGPUStoreOp_Discard };
WGPURenderPassColorAttachment colorAttachment{};
colorAttachment.view = target->texView,
colorAttachment.loadOp = clear ? WGPULoadOp_Clear : WGPULoadOp_Load,
colorAttachment.storeOp = WGPUStoreOp_Store;
#ifdef __EMSCRIPTEN__
colorAttachment.depthSlice = WGPU_DEPTH_SLICE_UNDEFINED;
#endif
WGPURenderPassDescriptor renderPassDesc{ .colorAttachmentCount = 1, .colorAttachments = &colorAttachment, .depthStencilAttachment = &depthStencilAttachment };
renderPassEncoder = wgpuCommandEncoderBeginRenderPass(commandEncoder, &renderPassDesc);
assert(renderPassEncoder);
}
void WgCompositor::endRenderPass()
{
assert(renderPassEncoder);
wgpuRenderPassEncoderEnd(renderPassEncoder);
wgpuRenderPassEncoderRelease(renderPassEncoder);
this->renderPassEncoder = nullptr;
this->currentTarget = nullptr;
}
void WgCompositor::renderClipPath(WgContext& context, WgRenderDataPaint* renderData, WgRenderStorage* dst)
{
assert(renderData);
if (renderData->clips.count == 0) return;
// store current render pass
WgRenderStorage *target = currentTarget;
endRenderPass();
// render first clip path
beginRenderPass(commandEncoder, dst, true);
drawClipPath(context, (WgRenderDataShape*)(renderData->clips[0]));
endRenderPass();
// render amd merge clip paths
for (uint32_t i = 1 ; i < renderData->clips.count; i++) {
// render clip path
beginRenderPass(commandEncoder, &storageInterm, true);
drawClipPath(context, (WgRenderDataShape*)(renderData->clips[i]));
endRenderPass();
// merge masks
mergeMasks(commandEncoder, &storageInterm, dst);
}
// restore current render pass
beginRenderPass(commandEncoder, target, false);
}
void WgCompositor::renderShape(WgContext& context, WgRenderDataShape* renderData, BlendMethod blendMethod)
{
assert(renderData);
assert(renderPassEncoder);
WgPipelineBlendType blentType = blendMethodToBlendType(blendMethod);
// apply clip path if neccessary
if (renderData->clips.count != 0) {
renderClipPath(context, renderData, &storageClipPath);
if (renderData->strokeFirst) {
composeStrokes(context, renderData, &storageClipPath, CompositeMethod::ClipPath);
composeShape(context, renderData, &storageClipPath, CompositeMethod::ClipPath);
} else {
composeShape(context, renderData, &storageClipPath, CompositeMethod::ClipPath);
composeStrokes(context, renderData, &storageClipPath, CompositeMethod::ClipPath);
}
// use custom blending
} else if (blentType == WgPipelineBlendType::Custom) {
if (renderData->strokeFirst) {
blendStrokes(context, renderData, blendMethod);
blendShape(context, renderData, blendMethod);
} else {
blendShape(context, renderData, blendMethod);
blendStrokes(context, renderData, blendMethod);
}
// use direct hardware blending
} else {
if (renderData->strokeFirst) {
drawStrokes(context, renderData, blentType);
drawShape(context, renderData, blentType);
} else {
drawShape(context, renderData, blentType);
drawStrokes(context, renderData, blentType);
}
}
}
void WgCompositor::renderImage(WgContext& context, WgRenderDataPicture* renderData, BlendMethod blendMethod)
{
assert(renderData);
assert(renderPassEncoder);
WgPipelineBlendType blentType = blendMethodToBlendType(blendMethod);
// apply clip path if neccessary
if (renderData->clips.count != 0) {
renderClipPath(context, renderData, &storageClipPath);
composeImage(context, renderData, &storageClipPath, CompositeMethod::ClipPath);
// use custom blending
} else if (blentType == WgPipelineBlendType::Custom)
blendImage(context, renderData, blendMethod);
else // use direct hardware blending
drawImage(context, renderData, blentType);
};
void WgCompositor::blendShape(WgContext& context, WgRenderDataShape* renderData, BlendMethod blendMethod)
{
assert(renderData);
assert(commandEncoder);
assert(currentTarget);
// skip shape blending if shapes do not exist
if (renderData->renderSettingsShape.skip) return;
if (renderData->meshGroupShapes.meshes.count == 0) return;
// store current render pass
WgRenderStorage *target = currentTarget;
endRenderPass();
// render into intermediate buffer
beginRenderPass(commandEncoder, &storageInterm, true);
drawShape(context, renderData, WgPipelineBlendType::Custom);
endRenderPass();
// run blend
blend(commandEncoder, &storageInterm, target, 255, blendMethod, renderData->renderSettingsShape.rasterType);
// restore current render pass
beginRenderPass(commandEncoder, target, false);
}
void WgCompositor::blendStrokes(WgContext& context, WgRenderDataShape* renderData, BlendMethod blendMethod)
{
assert(renderData);
assert(commandEncoder);
assert(currentTarget);
// skip strokes blending if strokes do not exist
if (renderData->renderSettingsStroke.skip) return;
if (renderData->meshGroupStrokes.meshes.count == 0) return;
// store current render pass
WgRenderStorage *target = currentTarget;
endRenderPass();
// render into intermediate buffer
beginRenderPass(commandEncoder, &storageInterm, true);
drawStrokes(context, renderData, WgPipelineBlendType::Custom);
endRenderPass();
// run blend
blend(commandEncoder, &storageInterm, target, 255, blendMethod, renderData->renderSettingsStroke.rasterType);
// restore current render pass
beginRenderPass(commandEncoder, target, false);
};
void WgCompositor::blendImage(WgContext& context, WgRenderDataPicture* renderData, BlendMethod blendMethod)
{
assert(renderData);
assert(commandEncoder);
assert(currentTarget);
// store current render pass
WgRenderStorage *target = currentTarget;
endRenderPass();
// render into intermediate buffer
beginRenderPass(commandEncoder, &storageInterm, true);
drawImage(context, renderData, WgPipelineBlendType::Custom);
endRenderPass();
// run blend
blend(commandEncoder, &storageInterm, target, 255, blendMethod, WgRenderRasterType::Image);
// restore current render pass
beginRenderPass(commandEncoder, target, false);
};
void WgCompositor::composeShape(WgContext& context, WgRenderDataShape* renderData, WgRenderStorage* mask, CompositeMethod composeMethod)
{
assert(mask);
assert(renderData);
assert(commandEncoder);
assert(currentTarget);
// skip shape composing if shape do not exist
if (renderData->renderSettingsShape.skip) return;
if (renderData->meshGroupShapes.meshes.count == 0) return;
// store current render pass
WgRenderStorage *target = currentTarget;
endRenderPass();
// render into intermediate buffer
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);
}
void WgCompositor::composeStrokes(WgContext& context, WgRenderDataShape* renderData, WgRenderStorage* msk, CompositeMethod composeMethod)
{
assert(msk);
assert(renderData);
assert(commandEncoder);
assert(currentTarget);
// skip shape composing if strokes do not exist
if (renderData->renderSettingsStroke.skip) return;
if (renderData->meshGroupStrokes.meshes.count == 0) return;
// store current render pass
WgRenderStorage *target = currentTarget;
endRenderPass();
// render into intermediate buffer
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);
}
void WgCompositor::composeImage(WgContext& context, WgRenderDataPicture* renderData, WgRenderStorage* msk, CompositeMethod composeMethod)
{
assert(msk);
assert(renderData);
assert(commandEncoder);
assert(currentTarget);
// store current render pass
WgRenderStorage *target = currentTarget;
endRenderPass();
// render into intermediate buffer
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);
}
void WgCompositor::drawClipPath(WgContext& context, WgRenderDataShape* renderData)
{
assert(renderData);
assert(renderPassEncoder);
assert(renderData->meshGroupShapes.meshes.count == renderData->meshGroupShapesBBox.meshes.count);
if ((renderData->viewport.w <= 0) || (renderData->viewport.h <= 0)) return;
wgpuRenderPassEncoderSetScissorRect(renderPassEncoder, renderData->viewport.x, renderData->viewport.y, renderData->viewport.w, renderData->viewport.h);
// setup stencil rules
WGPURenderPipeline stencilPipeline = (renderData->fillRule == FillRule::Winding) ? pipelines->winding : pipelines->evenodd;
wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 0);
wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 0, bindGroupViewMat, 0, nullptr);
wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 1, renderData->bindGroupPaint, 0, nullptr);
wgpuRenderPassEncoderSetPipeline(renderPassEncoder, stencilPipeline);
// draw to stencil (first pass)
for (uint32_t i = 0; i < renderData->meshGroupShapes.meshes.count; i++)
renderData->meshGroupShapes.meshes[i]->drawFan(context, renderPassEncoder);
// setup fill rules
wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 0);
wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 0, bindGroupViewMat, 0, nullptr);
wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 1, renderData->bindGroupPaint, 0, nullptr);
wgpuRenderPassEncoderSetPipeline(renderPassEncoder, pipelines->clipPath);
// draw to color (second pass)
renderData->meshDataBBox.drawFan(context, renderPassEncoder);
}
void WgCompositor::drawShape(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType)
{
assert(renderData);
assert(renderPassEncoder);
assert(renderData->meshGroupShapes.meshes.count == renderData->meshGroupShapesBBox.meshes.count);
uint32_t blendTypeInd = (uint32_t)blendType;
if (renderData->renderSettingsShape.skip) return;
if ((renderData->viewport.w <= 0) || (renderData->viewport.h <= 0)) return;
wgpuRenderPassEncoderSetScissorRect(renderPassEncoder, renderData->viewport.x, renderData->viewport.y, renderData->viewport.w, renderData->viewport.h);
// setup stencil rules
WGPURenderPipeline stencilPipeline = (renderData->fillRule == FillRule::Winding) ? pipelines->winding : pipelines->evenodd;
wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 0);
wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 0, bindGroupViewMat, 0, nullptr);
wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 1, renderData->bindGroupPaint, 0, nullptr);
wgpuRenderPassEncoderSetPipeline(renderPassEncoder, stencilPipeline);
// draw to stencil (first pass)
for (uint32_t i = 0; i < renderData->meshGroupShapes.meshes.count; i++)
renderData->meshGroupShapes.meshes[i]->drawFan(context, renderPassEncoder);
// setup fill rules
wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 0);
wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 0, bindGroupViewMat, 0, nullptr);
wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 1, renderData->bindGroupPaint, 0, nullptr);
WgRenderSettings& settings = renderData->renderSettingsShape;
if (settings.fillType == WgRenderSettingsType::Solid) {
wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 2, settings.bindGroupSolid, 0, nullptr);
wgpuRenderPassEncoderSetPipeline(renderPassEncoder, pipelines->solid[blendTypeInd]);
} else if (settings.fillType == WgRenderSettingsType::Linear) {
wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 2, settings.bindGroupLinear, 0, nullptr);
wgpuRenderPassEncoderSetPipeline(renderPassEncoder, pipelines->linear[blendTypeInd]);
} else if (settings.fillType == WgRenderSettingsType::Radial) {
wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 2, settings.bindGroupRadial, 0, nullptr);
wgpuRenderPassEncoderSetPipeline(renderPassEncoder, pipelines->radial[blendTypeInd]);
}
// draw to color (second pass)
renderData->meshDataBBox.drawFan(context, renderPassEncoder);
}
void WgCompositor::drawStrokes(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType)
{
assert(renderData);
assert(renderPassEncoder);
assert(renderData->meshGroupStrokes.meshes.count == renderData->meshGroupStrokesBBox.meshes.count);
uint32_t blendTypeInd = (uint32_t)blendType;
if (renderData->renderSettingsStroke.skip) return;
if ((renderData->viewport.w <= 0) || (renderData->viewport.h <= 0)) return;
wgpuRenderPassEncoderSetScissorRect(renderPassEncoder, renderData->viewport.x, renderData->viewport.y, renderData->viewport.w, renderData->viewport.h);
// draw strokes to stencil (first pass)
for (uint32_t i = 0; i < renderData->meshGroupStrokes.meshes.count; i++) {
// setup stencil rules
wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 255);
wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 0, bindGroupViewMat, 0, nullptr);
wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 1, renderData->bindGroupPaint, 0, nullptr);
wgpuRenderPassEncoderSetPipeline(renderPassEncoder, pipelines->direct);
// draw to stencil (first pass)
renderData->meshGroupStrokes.meshes[i]->draw(context, renderPassEncoder);
// setup fill rules
wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 0);
wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 0, bindGroupViewMat, 0, nullptr);
wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 1, renderData->bindGroupPaint, 0, nullptr);
WgRenderSettings& settings = renderData->renderSettingsStroke;
if (settings.fillType == WgRenderSettingsType::Solid) {
wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 2, settings.bindGroupSolid, 0, nullptr);
wgpuRenderPassEncoderSetPipeline(renderPassEncoder, pipelines->solid[blendTypeInd]);
} else if (settings.fillType == WgRenderSettingsType::Linear) {
wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 2, settings.bindGroupLinear, 0, nullptr);
wgpuRenderPassEncoderSetPipeline(renderPassEncoder, pipelines->linear[blendTypeInd]);
} else if (settings.fillType == WgRenderSettingsType::Radial) {
wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 2, settings.bindGroupRadial, 0, nullptr);
wgpuRenderPassEncoderSetPipeline(renderPassEncoder, pipelines->radial[blendTypeInd]);
}
// draw to color (second pass)
renderData->meshGroupStrokesBBox.meshes[i]->drawFan(context, renderPassEncoder);
}
}
void WgCompositor::drawImage(WgContext& context, WgRenderDataPicture* renderData, WgPipelineBlendType blendType)
{
assert(renderData);
assert(renderPassEncoder);
uint32_t blendTypeInd = (uint32_t)blendType;
if ((renderData->viewport.w <= 0) || (renderData->viewport.h <= 0)) return;
wgpuRenderPassEncoderSetScissorRect(renderPassEncoder, renderData->viewport.x, renderData->viewport.y, renderData->viewport.w, renderData->viewport.h);
wgpuRenderPassEncoderSetStencilReference(renderPassEncoder, 0);
wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 0, bindGroupViewMat, 0, nullptr);
wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 1, renderData->bindGroupPaint, 0, nullptr);
wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 2, renderData->bindGroupPicture, 0, nullptr);
wgpuRenderPassEncoderSetPipeline(renderPassEncoder, pipelines->image[blendTypeInd]);
renderData->meshData.drawImage(context, renderPassEncoder);
}
void WgCompositor::mergeMasks(WGPUCommandEncoder encoder, WgRenderStorage* mask0, WgRenderStorage* mask1)
{
assert(mask0);
assert(mask1);
assert(!renderPassEncoder);
// copy dst storage to temporary read only storage
const WGPUImageCopyTexture texSrc { .texture = mask1->texture };
const WGPUImageCopyTexture texDst { .texture = storageDstCopy.texture };
const WGPUExtent3D copySize { .width = width, .height = height, .depthOrArrayLayers = 1 };
wgpuCommandEncoderCopyTextureToTexture(encoder, &texSrc, &texDst, &copySize);
// execute compose shader
const WGPUComputePassDescriptor computePassDescriptor{};
WGPUComputePassEncoder computePassEncoder = wgpuCommandEncoderBeginComputePass(encoder, &computePassDescriptor);
wgpuComputePassEncoderSetBindGroup(computePassEncoder, 0, mask0->bindGroupRead, 0, nullptr);
wgpuComputePassEncoderSetBindGroup(computePassEncoder, 1, storageDstCopy.bindGroupRead, 0, nullptr);
wgpuComputePassEncoderSetBindGroup(computePassEncoder, 2, mask1->bindGroupWrite, 0, nullptr);
wgpuComputePassEncoderSetPipeline(computePassEncoder, pipelines->mergeMasks);
wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, (width + 7) / 8, (height + 7) / 8, 1);
wgpuComputePassEncoderEnd(computePassEncoder);
}
void WgCompositor::blend(WGPUCommandEncoder encoder, WgRenderStorage* src, WgRenderStorage* dst, uint8_t opacity, BlendMethod blendMethod, WgRenderRasterType rasterType)
{
assert(src);
assert(dst);
assert(!renderPassEncoder);
WGPUComputePipeline pipeline = pipelines->blendImage[(size_t)blendMethod];
if (rasterType == WgRenderRasterType::Solid) pipeline = pipelines->blendSolid[(size_t)blendMethod];
if (rasterType == WgRenderRasterType::Gradient) pipeline = pipelines->blendGradient[(size_t)blendMethod];
// copy texture to texture
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, &copySize);
// execute compose shader
const WGPUComputePassDescriptor computePassDescriptor{};
WGPUComputePassEncoder computePassEncoder = wgpuCommandEncoderBeginComputePass(encoder, &computePassDescriptor);
wgpuComputePassEncoderSetBindGroup(computePassEncoder, 0, src->bindGroupRead, 0, nullptr);
wgpuComputePassEncoderSetBindGroup(computePassEncoder, 1, storageDstCopy.bindGroupRead, 0, nullptr);
wgpuComputePassEncoderSetBindGroup(computePassEncoder, 2, dst->bindGroupWrite, 0, nullptr);
wgpuComputePassEncoderSetBindGroup(computePassEncoder, 3, bindGroupOpacities[opacity], 0, nullptr);
wgpuComputePassEncoderSetPipeline(computePassEncoder, pipeline);
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, &copySize);
// 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);
}

90
src/renderer/wg_engine/tvgWgCompositor.h Executable file
View file

@ -0,0 +1,90 @@
/*
* Copyright (c) 2024 the ThorVG project. All rights reserved.
* Permission is hereby granted, free of charge, to any person obtaining a copy
* of this software and associated documentation files (the "Software"), to deal
* in the Software without restriction, including without limitation the rights
* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
* copies of the Software, and to permit persons to whom the Software is
* furnished to do so, subject to the following conditions:
* The above copyright notice and this permission notice shall be included in all
* copies or substantial portions of the Software.
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
* SOFTWARE.
*/
#ifndef _TVG_WG_COMPOSITOR_H_
#define _TVG_WG_COMPOSITOR_H_
#include "tvgWgRenderTarget.h"
#include "tvgWgRenderData.h"
struct WgCompose: public Compositor {
BlendMethod blend;
};
class WgCompositor {
private:
// pipelines (external handle, do not release)
WgPipelines* pipelines{};
// global stencil buffer handles
WGPUTexture texStencil{};
WGPUTextureView texViewStencil{};
// global view matrix handles
WGPUBuffer bufferViewMat{};
WGPUBindGroup bindGroupViewMat{};
// opacity value pool
WGPUBuffer bufferOpacities[256]{};
WGPUBindGroup bindGroupOpacities[256]{};
// current render pass handles
WGPURenderPassEncoder renderPassEncoder{};
WGPUCommandEncoder commandEncoder{};
WgRenderStorage* currentTarget{};
// intermediate render storages
WgRenderStorage storageInterm;
WgRenderStorage storageClipPath;
WgRenderStorage storageDstCopy;
static WgPipelineBlendType blendMethodToBlendType(BlendMethod blendMethod);
public:
// render target dimensions
uint32_t width{};
uint32_t height{};
public:
void initialize(WgContext& context, uint32_t width, uint32_t height);
void release(WgContext& context);
void beginRenderPass(WGPUCommandEncoder encoder, WgRenderStorage* target, bool clear);
void endRenderPass();
void renderClipPath(WgContext& context, WgRenderDataPaint* renderData, WgRenderStorage* dst);
void renderShape(WgContext& context, WgRenderDataShape* renderData, BlendMethod blendMethod);
void renderImage(WgContext& context, WgRenderDataPicture* renderData, BlendMethod blendMethod);
void blendShape(WgContext& context, WgRenderDataShape* renderData, BlendMethod blendMethod);
void blendStrokes(WgContext& context, WgRenderDataShape* renderData, BlendMethod blendMethod);
void blendImage(WgContext& context, WgRenderDataPicture* renderData, BlendMethod blendMethod);
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 drawClipPath(WgContext& context, WgRenderDataShape* renderData);
void drawShape(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType);
void drawStrokes(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType);
void drawImage(WgContext& context, WgRenderDataPicture* renderData, WgPipelineBlendType blendType);
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_

0
src/renderer/wg_engine/tvgWgGeometry.cpp Normal file → Executable file
View file

0
src/renderer/wg_engine/tvgWgGeometry.h Normal file → Executable file
View file

747
src/renderer/wg_engine/tvgWgPipelines.cpp Normal file → Executable file
View file

@ -23,492 +23,359 @@
#include "tvgWgPipelines.h"
#include "tvgWgShaderSrc.h"
#define ARRAY_ELEMENTS_COUNT(arr) sizeof(arr)/sizeof(arr[0])
//************************************************************************
// graphics pipelines
//************************************************************************
void WgPipelineFillShapeWinding::initialize(WGPUDevice device)
WGPUShaderModule WgPipelines::createShaderModule(WGPUDevice device, const char* label, const char* code)
{
// vertex attributes settings
WGPUVertexAttribute vertexAttributesPos = { WGPUVertexFormat_Float32x2, sizeof(float) * 0, 0 };
WGPUVertexBufferLayout vertexBufferLayouts[] = {
makeVertexBufferLayout(&vertexAttributesPos, 1, sizeof(float) * 2)
};
// bind groups
WGPUBindGroupLayout bindGroupLayouts[] = {
WgBindGroupCanvas::getLayout(device),
WgBindGroupPaint::getLayout(device)
};
// stencil function
WGPUCompareFunction stencilFunctionFront = WGPUCompareFunction_Always;
WGPUStencilOperation stencilOperationFront = WGPUStencilOperation_IncrementWrap;
WGPUCompareFunction stencilFunctionBack = WGPUCompareFunction_Always;
WGPUStencilOperation stencilOperationBack = WGPUStencilOperation_DecrementWrap;
// shader source and labels
auto shaderSource = cShaderSource_PipelineFill;
auto shaderLabel = "The shader fill";
auto pipelineLabel = "The render pipeline fill shape winding";
// allocate all pipeline handles
allocate(device, WgPipelineBlendType::SrcOver, WGPUColorWriteMask_None,
vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts),
bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts),
stencilFunctionFront, stencilOperationFront, stencilFunctionBack, stencilOperationBack,
shaderSource, shaderLabel, pipelineLabel);
WGPUShaderModuleWGSLDescriptor shaderModuleWGSLDesc{};
shaderModuleWGSLDesc.chain.sType = WGPUSType_ShaderModuleWGSLDescriptor;
shaderModuleWGSLDesc.code = code;
const WGPUShaderModuleDescriptor shaderModuleDesc { .nextInChain = &shaderModuleWGSLDesc.chain, .label = label };
return wgpuDeviceCreateShaderModule(device, &shaderModuleDesc);
}
void WgPipelineFillShapeEvenOdd::initialize(WGPUDevice device)
WGPUPipelineLayout WgPipelines::createPipelineLayout(WGPUDevice device, const WGPUBindGroupLayout* bindGroupLayouts, const uint32_t bindGroupLayoutsCount)
{
// vertex attributes settings
WGPUVertexAttribute vertexAttributesPos = { WGPUVertexFormat_Float32x2, sizeof(float) * 0, 0 };
WGPUVertexBufferLayout vertexBufferLayouts[] = {
makeVertexBufferLayout(&vertexAttributesPos, 1, sizeof(float) * 2)
};
// bind groups
WGPUBindGroupLayout bindGroupLayouts[] = {
WgBindGroupCanvas::getLayout(device),
WgBindGroupPaint::getLayout(device)
};
// stencil function
WGPUCompareFunction stencilFunctionFront = WGPUCompareFunction_Always;
WGPUStencilOperation stencilOperationFront = WGPUStencilOperation_Invert;
WGPUCompareFunction stencilFunctionBack = WGPUCompareFunction_Always;
WGPUStencilOperation stencilOperationBack = WGPUStencilOperation_Invert;
// shader source and labels
auto shaderSource = cShaderSource_PipelineFill;
auto shaderLabel = "The shader fill";
auto pipelineLabel = "The render pipeline fill shape Even Odd";
// allocate all pipeline handles
allocate(device, WgPipelineBlendType::SrcOver, WGPUColorWriteMask_None,
vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts),
bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts),
stencilFunctionFront, stencilOperationFront, stencilFunctionBack, stencilOperationBack,
shaderSource, shaderLabel, pipelineLabel);
const WGPUPipelineLayoutDescriptor pipelineLayoutDesc { .bindGroupLayoutCount = bindGroupLayoutsCount, .bindGroupLayouts = bindGroupLayouts };
return wgpuDeviceCreatePipelineLayout(device, &pipelineLayoutDesc);
}
void WgPipelineFillStroke::initialize(WGPUDevice device)
WGPURenderPipeline WgPipelines::createRenderPipeline(
WGPUDevice device, const char* pipelineLabel,
const WGPUShaderModule shaderModule, const WGPUPipelineLayout pipelineLayout,
const WGPUVertexBufferLayout *vertexBufferLayouts, const uint32_t vertexBufferLayoutsCount,
const WGPUColorWriteMaskFlags writeMask,
const WGPUCompareFunction stencilFunctionFrnt, const WGPUStencilOperation stencilOperationFrnt,
const WGPUCompareFunction stencilFunctionBack, const WGPUStencilOperation stencilOperationBack,
const WGPUPrimitiveState primitiveState, const WGPUMultisampleState multisampleState, const WGPUBlendState blendState)
{
// vertex and buffers settings
WGPUVertexAttribute vertexAttributesPos = { WGPUVertexFormat_Float32x2, sizeof(float) * 0, 0 };
WGPUVertexBufferLayout vertexBufferLayouts[] = {
makeVertexBufferLayout(&vertexAttributesPos, 1, sizeof(float) * 2)
const WGPUColorTargetState colorTargetState { .format = WGPUTextureFormat_RGBA8Unorm, .blend = &blendState, .writeMask = writeMask };
const WGPUColorTargetState colorTargetStates[] { colorTargetState };
const WGPUDepthStencilState depthStencilState {
.format = WGPUTextureFormat_Stencil8, .depthCompare = WGPUCompareFunction_Always,
.stencilFront = { .compare = stencilFunctionFrnt, .failOp = stencilOperationFrnt, .depthFailOp = stencilOperationFrnt, .passOp = stencilOperationFrnt },
.stencilBack = { .compare = stencilFunctionBack, .failOp = stencilOperationBack, .depthFailOp = stencilOperationBack, .passOp = stencilOperationBack },
.stencilReadMask = 0xFFFFFFFF, .stencilWriteMask = 0xFFFFFFFF
};
// bind groups and layouts
WGPUBindGroupLayout bindGroupLayouts[] = {
WgBindGroupCanvas::getLayout(device),
WgBindGroupPaint::getLayout(device)
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 WGPURenderPipelineDescriptor renderPipelineDesc {
.label = pipelineLabel,
.layout = pipelineLayout,
.vertex = vertexState,
.primitive = primitiveState,
.depthStencil = &depthStencilState,
.multisample = multisampleState,
.fragment = &fragmentState
};
// stencil function
WGPUCompareFunction stencilFunction = WGPUCompareFunction_Always;
WGPUStencilOperation stencilOperation = WGPUStencilOperation_Replace;
// shader source and labels
auto shaderSource = cShaderSource_PipelineFill;
auto shaderLabel = "The shader fill";
auto pipelineLabel = "The render pipeline fill stroke";
// allocate all pipeline handles
allocate(device, WgPipelineBlendType::SrcOver, WGPUColorWriteMask_None,
vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts),
bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts),
stencilFunction, stencilOperation, stencilFunction, stencilOperation,
shaderSource, shaderLabel, pipelineLabel);
return wgpuDeviceCreateRenderPipeline(device, &renderPipelineDesc);
}
void WgPipelineClipMask::initialize(WGPUDevice device)
WGPUComputePipeline WgPipelines::createComputePipeline(
WGPUDevice device, const char* pipelineLabel,
const WGPUShaderModule shaderModule, const WGPUPipelineLayout pipelineLayout)
{
// vertex and buffers settings
WGPUVertexAttribute vertexAttributesPos = { WGPUVertexFormat_Float32x2, sizeof(float) * 0, 0 };
WGPUVertexBufferLayout vertexBufferLayouts[] = {
makeVertexBufferLayout(&vertexAttributesPos, 1, sizeof(float) * 2)
const WGPUComputePipelineDescriptor computePipelineDesc{
.label = pipelineLabel,
.layout = pipelineLayout,
.compute = {
.module = shaderModule,
.entryPoint = "cs_main"
}
};
// bind groups
WGPUBindGroupLayout bindGroupLayouts[] = {
WgBindGroupCanvas::getLayout(device),
WgBindGroupPaint::getLayout(device)
};
// stencil function
WGPUCompareFunction stencilFunction = WGPUCompareFunction_NotEqual;
WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero;
// shader source and labels
auto shaderSource = cShaderSource_PipelineFill;
auto shaderLabel = "The shader fill";
auto pipelineLabel = "The render pipeline clip mask";
// allocate all pipeline handles
allocate(device, WgPipelineBlendType::SrcOver, WGPUColorWriteMask_All,
vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts),
bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts),
stencilFunction, stencilOperation, stencilFunction, stencilOperation,
shaderSource, shaderLabel, pipelineLabel);
return wgpuDeviceCreateComputePipeline(device, &computePipelineDesc);
}
void WgPipelineSolid::initialize(WGPUDevice device, WgPipelineBlendType blendType)
void WgPipelines::releaseComputePipeline(WGPUComputePipeline& computePipeline)
{
// vertex and buffers settings
WGPUVertexAttribute vertexAttributesPos = { WGPUVertexFormat_Float32x2, sizeof(float) * 0, 0 };
WGPUVertexBufferLayout vertexBufferLayouts[] = {
makeVertexBufferLayout(&vertexAttributesPos, 1, sizeof(float) * 2)
};
// bind groups and layouts
WGPUBindGroupLayout bindGroupLayouts[] = {
WgBindGroupCanvas::getLayout(device),
WgBindGroupPaint::getLayout(device),
WgBindGroupSolidColor::getLayout(device)
};
// stencil function
WGPUCompareFunction stencilFunction = WGPUCompareFunction_NotEqual;
WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero;
// shader source and labels
auto shaderSource = cShaderSource_PipelineSolid;
auto shaderLabel = "The shader solid color";
auto pipelineLabel = "The render pipeline solid color";
// allocate all pipeline handles
allocate(device, blendType, WGPUColorWriteMask_All,
vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts),
bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts),
stencilFunction, stencilOperation, stencilFunction, stencilOperation,
shaderSource, shaderLabel, pipelineLabel);
if (computePipeline) {
wgpuComputePipelineRelease(computePipeline);
computePipeline = nullptr;
}
}
void WgPipelineLinear::initialize(WGPUDevice device, WgPipelineBlendType blendType)
void WgPipelines::releaseRenderPipeline(WGPURenderPipeline& renderPipeline)
{
// vertex and buffers settings
WGPUVertexAttribute vertexAttributesPos = { WGPUVertexFormat_Float32x2, sizeof(float) * 0, 0 };
WGPUVertexBufferLayout vertexBufferLayouts[] = {
makeVertexBufferLayout(&vertexAttributesPos, 1, sizeof(float) * 2)
};
// bind groups and layouts
WGPUBindGroupLayout bindGroupLayouts[] = {
WgBindGroupCanvas::getLayout(device),
WgBindGroupPaint::getLayout(device),
WgBindGroupLinearGradient::getLayout(device)
};
// stencil function
WGPUCompareFunction stencilFunction = WGPUCompareFunction_NotEqual;
WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero;
// shader source and labels
auto shaderSource = cShaderSource_PipelineLinear;
auto shaderLabel = "The shader linear gradient";
auto pipelineLabel = "The render pipeline linear gradient";
// allocate all pipeline handles
allocate(device, blendType, WGPUColorWriteMask_All,
vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts),
bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts),
stencilFunction, stencilOperation, stencilFunction, stencilOperation,
shaderSource, shaderLabel, pipelineLabel);
if (renderPipeline) {
wgpuRenderPipelineRelease(renderPipeline);
renderPipeline = nullptr;
}
}
void WgPipelineRadial::initialize(WGPUDevice device, WgPipelineBlendType blendType)
void WgPipelines::releasePipelineLayout(WGPUPipelineLayout& pipelineLayout)
{
// vertex and buffers settings
WGPUVertexAttribute vertexAttributesPos = { WGPUVertexFormat_Float32x2, sizeof(float) * 0, 0 };
WGPUVertexBufferLayout vertexBufferLayouts[] = {
makeVertexBufferLayout(&vertexAttributesPos, 1, sizeof(float) * 2)
};
// bind groups and layouts
WGPUBindGroupLayout bindGroupLayouts[] = {
WgBindGroupCanvas::getLayout(device),
WgBindGroupPaint::getLayout(device),
WgBindGroupRadialGradient::getLayout(device)
};
// stencil function
WGPUCompareFunction stencilFunction = WGPUCompareFunction_NotEqual;
WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero;
// shader source and labels
auto shaderSource = cShaderSource_PipelineRadial;
auto shaderLabel = "The shader radial gradient";
auto pipelineLabel = "The render pipeline radial gradient";
// allocate all pipeline handles
allocate(device, blendType, WGPUColorWriteMask_All,
vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts),
bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts),
stencilFunction, stencilOperation, stencilFunction, stencilOperation,
shaderSource, shaderLabel, pipelineLabel);
if (pipelineLayout) {
wgpuPipelineLayoutRelease(pipelineLayout);
pipelineLayout = nullptr;
}
}
void WgPipelineImage::initialize(WGPUDevice device, WgPipelineBlendType blendType)
void WgPipelines::releaseShaderModule(WGPUShaderModule& shaderModule)
{
// 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[] = {
WgBindGroupCanvas::getLayout(device),
WgBindGroupPaint::getLayout(device),
WgBindGroupPicture::getLayout(device)
};
// stencil function
WGPUCompareFunction stencilFunction = WGPUCompareFunction_Always;
WGPUStencilOperation stencilOperation = WGPUStencilOperation_Zero;
// shader source and labels
auto shaderSource = cShaderSource_PipelineImage;
auto shaderLabel = "The shader image";
auto pipelineLabel = "The render pipeline image";
// allocate all pipeline handles
allocate(device, blendType, WGPUColorWriteMask_All,
vertexBufferLayouts, ARRAY_ELEMENTS_COUNT(vertexBufferLayouts),
bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts),
stencilFunction, stencilOperation, stencilFunction, stencilOperation,
shaderSource, shaderLabel, pipelineLabel);
if (shaderModule) {
wgpuShaderModuleRelease(shaderModule);
shaderModule = nullptr;
}
}
//************************************************************************
// compute pipelines
//************************************************************************
void WgPipelineCopy::initialize(WGPUDevice device)
{
// bind groups and layouts
WGPUBindGroupLayout bindGroupLayouts[] = {
WgBindGroupTextureStorageRgbaRO::getLayout(device),
WgBindGroupTextureStorageRgbaWO::getLayout(device)
};
// shader source and labels
auto shaderSource = cShaderSource_PipelineComputeCopy;
auto shaderLabel = "The compute shader copy";
auto pipelineLabel = "The compute pipeline copy";
// allocate all pipeline handles
allocate(device,
bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts),
shaderSource, shaderLabel, pipelineLabel);
}
void WgPipelineBlend::initialize(WGPUDevice device, const char *shaderSource)
{
// bind groups and layouts
WGPUBindGroupLayout bindGroupLayouts[] = {
WgBindGroupTexBlend::getLayout(device),
WgBindGroupBlendMethod::getLayout(device),
WgBindGroupOpacity::getLayout(device)
};
// shader source and labels
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 WgPipelineBlendMask::initialize(WGPUDevice device, const char *shaderSource)
{
// bind groups and layouts
WGPUBindGroupLayout bindGroupLayouts[] = {
WgBindGroupTexBlendMask::getLayout(device),
WgBindGroupBlendMethod::getLayout(device),
WgBindGroupOpacity::getLayout(device)
};
// shader source and labels
auto shaderLabel = "The compute shader blend mask";
auto pipelineLabel = "The compute pipeline blend mask";
// allocate all pipeline handles
allocate(device,
bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts),
shaderSource, shaderLabel, pipelineLabel);
}
void WgPipelineMaskCompose::initialize(WGPUDevice device)
{
// bind groups and layouts
WGPUBindGroupLayout bindGroupLayouts[] = {
WgBindGroupTexMaskCompose::getLayout(device),
};
// shader source and labels
auto shaderSource = cShaderSource_PipelineComputeMaskCompose;
auto shaderLabel = "The compute shader mask compose";
auto pipelineLabel = "The compute pipeline mask compose";
// allocate all pipeline handles
allocate(device,
bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts),
shaderSource, shaderLabel, pipelineLabel);
}
void WgPipelineCompose::initialize(WGPUDevice device)
{
// bind groups and layouts
WGPUBindGroupLayout bindGroupLayouts[] = {
WgBindGroupTexCompose::getLayout(device),
WgBindGroupCompositeMethod::getLayout(device),
WgBindGroupBlendMethod::getLayout(device),
WgBindGroupOpacity::getLayout(device)
};
// shader source and labels
auto shaderSource = cShaderSource_PipelineComputeCompose;
auto shaderLabel = "The compute shader compose blend";
auto pipelineLabel = "The compute pipeline compose blend";
// allocate all pipeline handles
allocate(device,
bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts),
shaderSource, shaderLabel, pipelineLabel);
}
void WgPipelineAntiAliasing::initialize(WGPUDevice device)
{
// bind groups and layouts
WGPUBindGroupLayout bindGroupLayouts[] = {
WgBindGroupTextureStorageRgbaRO::getLayout(device),
WgBindGroupTextureStorageBgraWO::getLayout(device)
};
// shader source and labels
auto shaderSource = cShaderSource_PipelineComputeAntiAlias;
auto shaderLabel = "The compute shader anti-aliasing";
auto pipelineLabel = "The compute pipeline anti-aliasing";
// allocate all pipeline handles
allocate(device,
bindGroupLayouts, ARRAY_ELEMENTS_COUNT(bindGroupLayouts),
shaderSource, shaderLabel, pipelineLabel);
}
//************************************************************************
// pipelines
//************************************************************************
void WgPipelines::initialize(WgContext& context)
{
// fill pipelines
fillShapeWinding.initialize(context.device);
fillShapeEvenOdd.initialize(context.device);
fillStroke.initialize(context.device);
clipMask.initialize(context.device);
for (uint8_t type = (uint8_t)WgPipelineBlendType::SrcOver; type <= (uint8_t)WgPipelineBlendType::Custom; type++) {
solid[type].initialize(context.device, (WgPipelineBlendType)type);
linear[type].initialize(context.device, (WgPipelineBlendType)type);
radial[type].initialize(context.device, (WgPipelineBlendType)type);
image[type].initialize(context.device, (WgPipelineBlendType)type);
}
// compute pipelines
computeCopy.initialize(context.device);
computeBlendSolid.initialize(context.device, cShaderSource_PipelineComputeBlendSolid);
computeBlendGradient.initialize(context.device, cShaderSource_PipelineComputeBlendGradient);
computeBlendImage.initialize(context.device, cShaderSource_PipelineComputeBlendImage);
computeBlendSolidMask.initialize(context.device, cShaderSource_PipelineComputeBlendSolidMask);
computeBlendGradientMask.initialize(context.device, cShaderSource_PipelineComputeBlendGradientMask);
computeBlendImageMask.initialize(context.device, cShaderSource_PipelineComputeBlendImageMask);
computeMaskCompose.initialize(context.device);
computeCompose.initialize(context.device);
computeAntiAliasing.initialize(context.device);
// store pipelines to context
// share pipelines to context
assert(!context.pipelines);
context.pipelines = this;
}
// initialize bind group layouts
layouts.initialize(context);
void WgPipelines::release()
{
WgBindGroupTexCompose::releaseLayout();
WgBindGroupTexMaskCompose::releaseLayout();
WgBindGroupTexBlendMask::releaseLayout();
WgBindGroupTexBlend::releaseLayout();
WgBindGroupTextureStorageBgraRO::releaseLayout();
WgBindGroupTextureStorageBgraWO::releaseLayout();
WgBindGroupTextureStorageRgbaRO::releaseLayout();
WgBindGroupTextureStorageRgbaWO::releaseLayout();
WgBindGroupTexture::releaseLayout();
WgBindGroupOpacity::releaseLayout();
WgBindGroupPicture::releaseLayout();
WgBindGroupRadialGradient::releaseLayout();
WgBindGroupLinearGradient::releaseLayout();
WgBindGroupSolidColor::releaseLayout();
WgBindGroupPaint::releaseLayout();
WgBindGroupCanvas::releaseLayout();
// compute pipelines
computeAntiAliasing.release();
computeCompose.release();
computeBlendImageMask.release();
computeBlendGradientMask.release();
computeBlendSolidMask.release();
computeBlendImage.release();
computeBlendGradient.release();
computeBlendSolid.release();
computeCopy.release();
// fill pipelines
for (uint8_t type = (uint8_t)WgPipelineBlendType::SrcOver; type <= (uint8_t)WgPipelineBlendType::Custom; type++) {
image[type].release();
radial[type].release();
linear[type].release();
solid[type].release();
// common pipeline settings
const WGPUVertexAttribute vertexAttributePos { .format = WGPUVertexFormat_Float32x2, .offset = 0, .shaderLocation = 0 };
const WGPUVertexAttribute vertexAttributeTex { .format = WGPUVertexFormat_Float32x2, .offset = 0, .shaderLocation = 1 };
const WGPUVertexAttribute vertexAttributesPos[] { vertexAttributePos };
const WGPUVertexAttribute vertexAttributesTex[] { vertexAttributeTex };
const WGPUVertexBufferLayout vertexBufferLayoutPos { .arrayStride = 8, .stepMode = WGPUVertexStepMode_Vertex, .attributeCount = 1, .attributes = vertexAttributesPos };
const WGPUVertexBufferLayout vertexBufferLayoutTex { .arrayStride = 8, .stepMode = WGPUVertexStepMode_Vertex, .attributeCount = 1, .attributes = vertexAttributesTex };
const WGPUVertexBufferLayout vertexBufferLayoutsShape[] { vertexBufferLayoutPos };
const WGPUVertexBufferLayout vertexBufferLayoutsImage[] { vertexBufferLayoutPos, vertexBufferLayoutTex };
const WGPUPrimitiveState primitiveState { .topology = WGPUPrimitiveTopology_TriangleList };
const WGPUMultisampleState multisampleState { .count = 1, .mask = 0xFFFFFFFF, .alphaToCoverageEnabled = false };
// blend states
const WGPUBlendState blendStateSrc {
.color = { .operation = WGPUBlendOperation_Add, .srcFactor = WGPUBlendFactor_One, .dstFactor = WGPUBlendFactor_Zero },
.alpha = { .operation = WGPUBlendOperation_Add, .srcFactor = WGPUBlendFactor_One, .dstFactor = WGPUBlendFactor_Zero }
};
const WGPUBlendState blendStateNrm {
.color = { .operation = WGPUBlendOperation_Add, .srcFactor = WGPUBlendFactor_One, .dstFactor = WGPUBlendFactor_OneMinusSrcAlpha },
.alpha = { .operation = WGPUBlendOperation_Add, .srcFactor = WGPUBlendFactor_One, .dstFactor = WGPUBlendFactor_OneMinusSrcAlpha }
};
const WGPUBlendState blendStates[] {
blendStateSrc, // WgPipelineBlendType::SrcOver
blendStateNrm, // WgPipelineBlendType::Normal
blendStateSrc // WgPipelineBlendType::Custom (same as SrcOver)
};
// bind group layouts
const WGPUBindGroupLayout bindGroupLayoutsStencil[] = {
layouts.layoutBuffer1Un,
layouts.layoutBuffer2Un
};
const WGPUBindGroupLayout bindGroupLayoutsFill[] = {
layouts.layoutBuffer1Un,
layouts.layoutBuffer2Un,
layouts.layoutBuffer1Un
};
const WGPUBindGroupLayout bindGroupLayoutsImage[] = {
layouts.layoutBuffer1Un,
layouts.layoutBuffer2Un,
layouts.layoutTexSampled
};
const WGPUBindGroupLayout bindGroupLayoutsMergeMasks[] = {
layouts.layoutTexStrorage1RO,
layouts.layoutTexStrorage1RO,
layouts.layoutTexStrorage1WO
};
const WGPUBindGroupLayout bindGroupLayoutsBlend[] = {
layouts.layoutTexStrorage1RO,
layouts.layoutTexStrorage1RO,
layouts.layoutTexStrorage1WO,
layouts.layoutBuffer1Un
};
const WGPUBindGroupLayout bindGroupLayoutsCompose[] = {
layouts.layoutTexStrorage1RO,
layouts.layoutTexStrorage1RO,
layouts.layoutTexStrorage1RO,
layouts.layoutTexStrorage1WO
};
const WGPUBindGroupLayout bindGroupLayoutsCopy[] = {
layouts.layoutTexStrorage1RO,
layouts.layoutTexScreen1WO
};
// pipeline layouts
layoutStencil = createPipelineLayout(context.device, bindGroupLayoutsStencil, 2);
layoutFill = createPipelineLayout(context.device, bindGroupLayoutsFill, 3);
layoutImage = createPipelineLayout(context.device, bindGroupLayoutsImage, 3);
layoutBlend = createPipelineLayout(context.device, bindGroupLayoutsBlend, 4);
layoutCompose = createPipelineLayout(context.device, bindGroupLayoutsCompose, 4);
layoutMergeMasks = createPipelineLayout(context.device, bindGroupLayoutsMergeMasks, 3);
layoutCopy = createPipelineLayout(context.device, bindGroupLayoutsCopy, 2);
// graphics shader modules
shaderStencil = createShaderModule(context.device, "The shader stencil", cShaderSrc_Stencil);
shaderSolid = createShaderModule(context.device, "The shader solid", cShaderSrc_Solid);
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);
// computes shader modules
shaderMergeMasks = createShaderModule(context.device, "The shader merge mask", cShaderSrc_MergeMasks);
shaderCopy = createShaderModule(context.device, "The shader copy", cShaderSrc_Copy);
// render pipeline winding
winding = createRenderPipeline(
context.device, "The render pipeline winding",
shaderStencil, layoutStencil,
vertexBufferLayoutsShape, 1,
WGPUColorWriteMask_None,
WGPUCompareFunction_Always, WGPUStencilOperation_IncrementWrap,
WGPUCompareFunction_Always, WGPUStencilOperation_DecrementWrap,
primitiveState, multisampleState, blendStateSrc);
// render pipeline even-odd
evenodd = createRenderPipeline(
context.device, "The render pipeline even-odd",
shaderStencil, layoutStencil,
vertexBufferLayoutsShape, 1,
WGPUColorWriteMask_None,
WGPUCompareFunction_Always, WGPUStencilOperation_Invert,
WGPUCompareFunction_Always, WGPUStencilOperation_Invert,
primitiveState, multisampleState, blendStateSrc);
// render pipeline direct
direct = createRenderPipeline(
context.device, "The render pipeline direct",
shaderStencil, layoutStencil,
vertexBufferLayoutsShape, 1,
WGPUColorWriteMask_None,
WGPUCompareFunction_Always, WGPUStencilOperation_Replace,
WGPUCompareFunction_Always, WGPUStencilOperation_Replace,
primitiveState, multisampleState, blendStateSrc);
// render pipeline clip path
clipPath = createRenderPipeline(
context.device, "The render pipeline clip path",
shaderStencil, layoutStencil,
vertexBufferLayoutsShape, 1,
WGPUColorWriteMask_All,
WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero,
WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero,
primitiveState, multisampleState, blendStateSrc);
// render pipeline solid
for (uint32_t i = 0; i < 3; i++) {
solid[i] = createRenderPipeline(
context.device, "The render pipeline solid",
shaderSolid, layoutFill,
vertexBufferLayoutsShape, 1,
WGPUColorWriteMask_All,
WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero,
WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero,
primitiveState, multisampleState, blendStates[i]);
}
// render pipeline radial
for (uint32_t i = 0; i < 3; i++) {
radial[i] = createRenderPipeline(
context.device, "The render pipeline radial",
shaderRadial, layoutFill,
vertexBufferLayoutsShape, 1,
WGPUColorWriteMask_All,
WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero,
WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero,
primitiveState, multisampleState, blendStates[i]);
}
// render pipeline linear
for (uint32_t i = 0; i < 3; i++) {
linear[i] = createRenderPipeline(
context.device, "The render pipeline linear",
shaderLinear, layoutFill,
vertexBufferLayoutsShape, 1,
WGPUColorWriteMask_All,
WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero,
WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero,
primitiveState, multisampleState, blendStates[i]);
}
// render pipeline image
for (uint32_t i = 0; i < 3; i++) {
image[i] = createRenderPipeline(
context.device, "The render pipeline image",
shaderImage, layoutImage,
vertexBufferLayoutsImage, 2,
WGPUColorWriteMask_All,
WGPUCompareFunction_Always, WGPUStencilOperation_Zero,
WGPUCompareFunction_Always, WGPUStencilOperation_Zero,
primitiveState, multisampleState, blendStates[i]);
}
// compute pipelines
mergeMasks = createComputePipeline(context.device, "The pipeline merge masks", shaderMergeMasks, layoutMergeMasks);
copy = createComputePipeline(context.device, "The pipeline copy", shaderCopy, layoutCopy);
// compute pipelines blend
const size_t shaderBlendCnt = sizeof(cShaderSrc_Blend_Solid)/sizeof(cShaderSrc_Blend_Solid[0]);
for (uint32_t i = 0; i < shaderBlendCnt; i++) {
// blend shaders
shaderBlendSolid[i] = createShaderModule(context.device, "The shader blend solid", cShaderSrc_Blend_Solid[i]);
shaderBlendGradient[i] = createShaderModule(context.device, "The shader blend gradient", cShaderSrc_Blend_Gradient[i]);
shaderBlendImage[i] = createShaderModule(context.device, "The shader blend image", cShaderSrc_Blend_Image[i]);
// blend pipelines
blendSolid[i] = createComputePipeline(context.device, "The pipeline blend solid", shaderBlendSolid[i], layoutBlend);
blendGradient[i] = createComputePipeline(context.device, "The pipeline blend gradient", shaderBlendGradient[i], layoutBlend);
blendImage[i] = createComputePipeline(context.device, "The pipeline blend image", shaderBlendImage[i], layoutBlend);
}
// compute pipelines compose
const size_t shaderComposeCnt = sizeof(cShaderSrc_Compose)/sizeof(cShaderSrc_Compose[0]);
for (uint32_t i = 0; i < shaderComposeCnt; i++) {
shaderCompose[i] = createShaderModule(context.device, "The shader compose", cShaderSrc_Compose[i]);
compose[i] = createComputePipeline(context.device, "The pipeline compose", shaderCompose[i], layoutCompose);
}
clipMask.release();
fillStroke.release();
fillShapeEvenOdd.release();
fillShapeWinding.release();
}
bool WgPipelines::isBlendMethodSupportsHW(BlendMethod blendMethod)
void WgPipelines::releaseGraphicHandles(WgContext& context)
{
switch (blendMethod) {
case BlendMethod::SrcOver:
case BlendMethod::Normal:
return true;
default: return false;
};
releaseRenderPipeline(clipPath);
for (uint32_t i = 0; i < 3; i++) {
releaseRenderPipeline(image[i]);
releaseRenderPipeline(linear[i]);
releaseRenderPipeline(radial[i]);
releaseRenderPipeline(solid[i]);
}
releaseRenderPipeline(direct);
releaseRenderPipeline(evenodd);
releaseRenderPipeline(winding);
releasePipelineLayout(layoutImage);
releasePipelineLayout(layoutFill);
releasePipelineLayout(layoutStencil);
releaseShaderModule(shaderImage);
releaseShaderModule(shaderLinear);
releaseShaderModule(shaderRadial);
releaseShaderModule(shaderSolid);
releaseShaderModule(shaderStencil);
}
WgPipelineBlendType WgPipelines::blendMethodToBlendType(BlendMethod blendMethod)
void WgPipelines::releaseComputeHandles(WgContext& context)
{
switch (blendMethod) {
case BlendMethod::SrcOver: return WgPipelineBlendType::SrcOver;
case BlendMethod::Normal: return WgPipelineBlendType::Normal;
default: return WgPipelineBlendType::Custom;
};
const size_t shaderComposeCnt = sizeof(shaderCompose)/sizeof(shaderCompose[0]);
for (uint32_t i = 0; i < shaderComposeCnt; i++) {
releaseComputePipeline(compose[i]);
releaseShaderModule(shaderCompose[i]);
}
const size_t shaderBlendImageCnt = sizeof(shaderBlendImage)/sizeof(shaderBlendImage[0]);
for (uint32_t i = 0; i < shaderBlendImageCnt; i++) {
releaseComputePipeline(blendImage[i]);
releaseComputePipeline(blendSolid[i]);
releaseComputePipeline(blendGradient[i]);
releaseShaderModule(shaderBlendImage[i]);
releaseShaderModule(shaderBlendGradient[i]);
releaseShaderModule(shaderBlendSolid[i]);
}
releaseComputePipeline(copy);
releaseComputePipeline(mergeMasks);
releasePipelineLayout(layoutCompose);
releasePipelineLayout(layoutBlend);
releasePipelineLayout(layoutMergeMasks);
releasePipelineLayout(layoutCopy);
releaseShaderModule(shaderMergeMasks);
releaseShaderModule(shaderCopy);
}
void WgPipelines::release(WgContext& context)
{
releaseComputeHandles(context);
releaseGraphicHandles(context);
layouts.release(context);
}

278
src/renderer/wg_engine/tvgWgPipelines.h Normal file → Executable file
View file

@ -25,219 +25,75 @@
#include "tvgWgBindGroups.h"
//*****************************************************************************
// render pipelines
//*****************************************************************************
struct WgPipelineFillShapeWinding: public WgRenderPipeline
{
void initialize(WGPUDevice device) override;
void use(WGPURenderPassEncoder encoder, WgBindGroupCanvas& groupCanvas, WgBindGroupPaint& groupPaint)
{
set(encoder);
groupCanvas.set(encoder, 0);
groupPaint.set(encoder, 1);
}
};
struct WgPipelineFillShapeEvenOdd: public WgRenderPipeline
{
void initialize(WGPUDevice device) override;
void use(WGPURenderPassEncoder encoder, WgBindGroupCanvas& groupCanvas, WgBindGroupPaint& groupPaint)
{
set(encoder);
groupCanvas.set(encoder, 0);
groupPaint.set(encoder, 1);
}
};
struct WgPipelineFillStroke: public WgRenderPipeline
{
void initialize(WGPUDevice device) override;
void use(WGPURenderPassEncoder encoder, WgBindGroupCanvas& groupCanvas, WgBindGroupPaint& groupPaint)
{
set(encoder);
groupCanvas.set(encoder, 0);
groupPaint.set(encoder, 1);
}
};
struct WgPipelineClipMask: public WgRenderPipeline
{
void initialize(WGPUDevice device) override;
void use(WGPURenderPassEncoder encoder, WgBindGroupCanvas& groupCanvas,WgBindGroupPaint& groupPaint)
{
set(encoder);
groupCanvas.set(encoder, 0);
groupPaint.set(encoder, 1);
}
};
struct WgPipelineSolid: public WgRenderPipeline
{
void initialize(WGPUDevice device) override {}
void initialize(WGPUDevice device, WgPipelineBlendType blendType);
void use(WGPURenderPassEncoder encoder, WgBindGroupCanvas& groupCanvas,WgBindGroupPaint& groupPaint, WgBindGroupSolidColor& groupSolid)
{
set(encoder);
groupCanvas.set(encoder, 0);
groupPaint.set(encoder, 1);
groupSolid.set(encoder, 2);
}
};
struct WgPipelineLinear: public WgRenderPipeline
{
void initialize(WGPUDevice device) override {}
void initialize(WGPUDevice device, WgPipelineBlendType blendType);
void use(WGPURenderPassEncoder encoder, WgBindGroupCanvas& groupCanvas, WgBindGroupPaint& groupPaint, WgBindGroupLinearGradient& groupLinear)
{
set(encoder);
groupCanvas.set(encoder, 0);
groupPaint.set(encoder, 1);
groupLinear.set(encoder, 2);
}
};
struct WgPipelineRadial: public WgRenderPipeline
{
void initialize(WGPUDevice device) override {}
void initialize(WGPUDevice device, WgPipelineBlendType blendType);
void use(WGPURenderPassEncoder encoder, WgBindGroupCanvas& groupCanvas, WgBindGroupPaint& groupPaint, WgBindGroupRadialGradient& groupRadial)
{
set(encoder);
groupCanvas.set(encoder, 0);
groupPaint.set(encoder, 1);
groupRadial.set(encoder, 2);
}
};
struct WgPipelineImage: public WgRenderPipeline
{
void initialize(WGPUDevice device) override { assert(false); };
void initialize(WGPUDevice device, WgPipelineBlendType blendType);
void use(WGPURenderPassEncoder encoder, WgBindGroupCanvas& groupCanvas, WgBindGroupPaint& groupPaint, WgBindGroupPicture& groupPicture)
{
set(encoder);
groupCanvas.set(encoder, 0);
groupPaint.set(encoder, 1);
groupPicture.set(encoder, 2);
}
};
//*****************************************************************************
// compute pipelines
//*****************************************************************************
struct WgPipelineCopy: public WgComputePipeline
{
void initialize(WGPUDevice device) override;
void use(WGPUComputePassEncoder encoder, WgBindGroupTextureStorageRgbaRO& groupTexSrc, WgBindGroupTextureStorageRgbaWO& groupTexDst)
{
set(encoder);
groupTexSrc.set(encoder, 0);
groupTexDst.set(encoder, 1);
}
};
struct WgPipelineBlend: public WgComputePipeline
{
void initialize(WGPUDevice device) override { assert(false); };
void initialize(WGPUDevice device, const char *shaderSource);
void use(WGPUComputePassEncoder encoder, WgBindGroupTexBlend& groupTexBlend, WgBindGroupBlendMethod& blendMethod, WgBindGroupOpacity& groupOpacity)
{
set(encoder);
groupTexBlend.set(encoder, 0);
blendMethod.set(encoder, 1);
groupOpacity.set(encoder, 2);
}
};
struct WgPipelineBlendMask: public WgComputePipeline
{
void initialize(WGPUDevice device) override { assert(false); };
void initialize(WGPUDevice device, const char *shaderSource);
void use(WGPUComputePassEncoder encoder, WgBindGroupTexBlendMask& groupTexBlendMask, WgBindGroupBlendMethod& blendMethod, WgBindGroupOpacity& groupOpacity)
{
set(encoder);
groupTexBlendMask.set(encoder, 0);
blendMethod.set(encoder, 1);
groupOpacity.set(encoder, 2);
}
};
struct WgPipelineMaskCompose: public WgComputePipeline
{
void initialize(WGPUDevice device) override;
void use(WGPUComputePassEncoder encoder, WgBindGroupTexMaskCompose& groupTexMaskCompose)
{
set(encoder);
groupTexMaskCompose.set(encoder, 0);
}
};
struct WgPipelineCompose: public WgComputePipeline
{
void initialize(WGPUDevice device) override;
void use(WGPUComputePassEncoder encoder, WgBindGroupTexCompose& groupTexCompose, WgBindGroupCompositeMethod& groupComposeMethod, WgBindGroupBlendMethod& groupBlendMethod, WgBindGroupOpacity& groupOpacity)
{
set(encoder);
groupTexCompose.set(encoder, 0);
groupComposeMethod.set(encoder, 1);
groupBlendMethod.set(encoder, 2);
groupOpacity.set(encoder, 3);
}
};
struct WgPipelineAntiAliasing: public WgComputePipeline
{
void initialize(WGPUDevice device) override;
void use(WGPUComputePassEncoder encoder, WgBindGroupTextureStorageRgbaRO& groupTexSrc, WgBindGroupTextureStorageBgraWO& groupTexDst)
{
set(encoder);
groupTexSrc.set(encoder, 0);
groupTexDst.set(encoder, 1);
}
};
//*****************************************************************************
// pipelines
//*****************************************************************************
struct WgPipelines
{
// render pipelines
WgPipelineFillShapeWinding fillShapeWinding;
WgPipelineFillShapeEvenOdd fillShapeEvenOdd;
WgPipelineFillStroke fillStroke;
// fill pipelines
WgPipelineClipMask clipMask;
WgPipelineSolid solid[3];
WgPipelineLinear linear[3];
WgPipelineRadial radial[3];
WgPipelineImage image[3];
// compute pipelines
WgPipelineCopy computeCopy;
WgPipelineBlend computeBlendSolid;
WgPipelineBlend computeBlendGradient;
WgPipelineBlend computeBlendImage;
WgPipelineBlendMask computeBlendSolidMask;
WgPipelineBlendMask computeBlendGradientMask;
WgPipelineBlendMask computeBlendImageMask;
WgPipelineMaskCompose computeMaskCompose;
WgPipelineCompose computeCompose;
WgPipelineAntiAliasing computeAntiAliasing;
enum class WgPipelineBlendType { SrcOver = 0, Normal, Custom };
class WgPipelines {
private:
// graphics pipeline shaders
WGPUShaderModule shaderStencil{};
WGPUShaderModule shaderSolid{};
WGPUShaderModule shaderRadial{};
WGPUShaderModule shaderLinear{};
WGPUShaderModule shaderImage{};
// compute pipeline shaders
WGPUShaderModule shaderMergeMasks;
WGPUShaderModule shaderBlendSolid[14];
WGPUShaderModule shaderBlendGradient[14];
WGPUShaderModule shaderBlendImage[14];
WGPUShaderModule shaderCompose[10];
WGPUShaderModule shaderCopy;
private:
// graphics pipeline layouts
WGPUPipelineLayout layoutStencil{};
WGPUPipelineLayout layoutFill{};
WGPUPipelineLayout layoutImage{};
// compute pipeline layouts
WGPUPipelineLayout layoutMergeMasks{};
WGPUPipelineLayout layoutBlend{};
WGPUPipelineLayout layoutCompose{};
WGPUPipelineLayout layoutCopy{};
public:
// graphics pipeline
WgBindGroupLayouts layouts;
WGPURenderPipeline winding{};
WGPURenderPipeline evenodd{};
WGPURenderPipeline direct{};
WGPURenderPipeline solid[3]{};
WGPURenderPipeline radial[3]{};
WGPURenderPipeline linear[3]{};
WGPURenderPipeline image[3]{};
WGPURenderPipeline clipPath{};
// compute pipeline
WGPUComputePipeline mergeMasks;
WGPUComputePipeline blendSolid[14];
WGPUComputePipeline blendGradient[14];
WGPUComputePipeline blendImage[14];
WGPUComputePipeline compose[10];
WGPUComputePipeline copy;
private:
void releaseGraphicHandles(WgContext& context);
void releaseComputeHandles(WgContext& context);
private:
WGPUShaderModule createShaderModule(WGPUDevice device, const char* label, const char* code);
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 WGPUVertexBufferLayout *vertexBufferLayouts, const uint32_t vertexBufferLayoutsCount,
const WGPUColorWriteMaskFlags writeMask,
const WGPUCompareFunction stencilFunctionFrnt, const WGPUStencilOperation stencilOperationFrnt,
const WGPUCompareFunction stencilFunctionBack, const WGPUStencilOperation stencilOperationBack,
const WGPUPrimitiveState primitiveState, const WGPUMultisampleState multisampleState, const WGPUBlendState blendState);
WGPUComputePipeline createComputePipeline(
WGPUDevice device, const char* pipelineLabel,
const WGPUShaderModule shaderModule, const WGPUPipelineLayout pipelineLayout);
void releaseComputePipeline(WGPUComputePipeline& computePipeline);
void releaseRenderPipeline(WGPURenderPipeline& renderPipeline);
void releasePipelineLayout(WGPUPipelineLayout& pipelineLayout);
void releaseShaderModule(WGPUShaderModule& shaderModule);
public:
void initialize(WgContext& context);
void release();
static bool isBlendMethodSupportsHW(BlendMethod blendMethod);
static WgPipelineBlendType blendMethodToBlendType(BlendMethod blendMethod);
void release(WgContext& context);
};
#endif // _TVG_WG_PIPELINES_H_

93
src/renderer/wg_engine/tvgWgRenderData.cpp Normal file → Executable file
View file

@ -21,11 +21,9 @@
* SOFTWARE.
*/
#ifndef _TVG_WG_RENDER_DATA_H_
#define _TVG_WG_RENDER_DATA_H_
#include <algorithm>
#include "tvgWgRenderData.h"
#include "tvgWgShaderTypes.h"
//***********************************************************************
// WgMeshData
@ -42,7 +40,7 @@ void WgMeshData::draw(WgContext& context, WGPURenderPassEncoder renderPassEncode
void WgMeshData::drawFan(WgContext& context, WGPURenderPassEncoder renderPassEncoder)
{
wgpuRenderPassEncoderSetVertexBuffer(renderPassEncoder, 0, bufferPosition, 0, vertexCount * sizeof(float) * 2);
wgpuRenderPassEncoderSetIndexBuffer(renderPassEncoder, context.indexBufferFan, WGPUIndexFormat_Uint32, 0, indexCount * sizeof(uint32_t));
wgpuRenderPassEncoderSetIndexBuffer(renderPassEncoder, context.bufferIndexFan, WGPUIndexFormat_Uint32, 0, indexCount * sizeof(uint32_t));
wgpuRenderPassEncoderDrawIndexed(renderPassEncoder, indexCount, 1, 0, 0, 0);
}
@ -62,8 +60,8 @@ void WgMeshData::update(WgContext& context, const WgPolyline* polyline)
assert(polyline->pts.count > 2);
vertexCount = polyline->pts.count;
indexCount = (polyline->pts.count - 2) * 3;
context.allocateVertexBuffer(bufferPosition, &polyline->pts[0], vertexCount * sizeof(float) * 2);
context.allocateIndexBufferFan(vertexCount);
context.allocateBufferVertex(bufferPosition, (float *)&polyline->pts[0], vertexCount * sizeof(float) * 2);
context.allocateBufferIndexFan(vertexCount);
}
@ -74,13 +72,13 @@ void WgMeshData::update(WgContext& context, const WgGeometryData* geometryData)
indexCount = geometryData->indexes.count;
// buffer position data create and write
if (geometryData->positions.pts.count > 0)
context.allocateVertexBuffer(bufferPosition, &geometryData->positions.pts[0], vertexCount * sizeof(float) * 2);
context.allocateBufferVertex(bufferPosition, (float *)&geometryData->positions.pts[0], vertexCount * sizeof(float) * 2);
// buffer tex coords data create and write
if (geometryData->texCoords.count > 0)
context.allocateVertexBuffer(bufferTexCoord, &geometryData->texCoords[0], vertexCount * sizeof(float) * 2);
context.allocateBufferVertex(bufferTexCoord, (float *)&geometryData->texCoords[0], vertexCount * sizeof(float) * 2);
// buffer index data create and write
if (geometryData->indexes.count > 0)
context.allocateIndexBuffer(bufferIndex, &geometryData->indexes[0], indexCount * sizeof(uint32_t));
context.allocateBufferIndex(bufferIndex, &geometryData->indexes[0], indexCount * sizeof(uint32_t));
};
@ -92,8 +90,8 @@ void WgMeshData::update(WgContext& context, const WgPoint pmin, const WgPoint pm
pmin.x, pmin.y, pmax.x, pmin.y,
pmax.x, pmax.y, pmin.x, pmax.y
};
context.allocateVertexBuffer(bufferPosition, data, sizeof(data));
context.allocateIndexBufferFan(vertexCount);
context.allocateBufferVertex(bufferPosition, data, sizeof(data));
context.allocateBufferIndexFan(vertexCount);
}
@ -186,29 +184,14 @@ void WgImageData::update(WgContext& context, Surface* surface)
{
release(context);
assert(surface);
texture = context.createTexture2d(
WGPUTextureUsage_TextureBinding | WGPUTextureUsage_CopyDst,
WGPUTextureFormat_RGBA8Unorm,
surface->w, surface->h, "The shape texture");
texture = context.createTexture(surface->w, surface->h, WGPUTextureFormat_RGBA8Unorm);
assert(texture);
textureView = context.createTextureView2d(texture, "The shape texture view");
textureView = context.createTextureView(texture);
assert(textureView);
// update texture data
WGPUImageCopyTexture imageCopyTexture{};
imageCopyTexture.nextInChain = nullptr;
imageCopyTexture.texture = texture;
imageCopyTexture.mipLevel = 0;
imageCopyTexture.origin = { 0, 0, 0 };
imageCopyTexture.aspect = WGPUTextureAspect_All;
WGPUTextureDataLayout textureDataLayout{};
textureDataLayout.nextInChain = nullptr;
textureDataLayout.offset = 0;
textureDataLayout.bytesPerRow = 4 * surface->w;
textureDataLayout.rowsPerImage = surface->h;
WGPUExtent3D writeSize{};
writeSize.width = surface->w;
writeSize.height = surface->h;
writeSize.depthOrArrayLayers = 1;
WGPUImageCopyTexture imageCopyTexture{ .texture = texture };
WGPUTextureDataLayout textureDataLayout{ .bytesPerRow = 4 * surface->w, .rowsPerImage = surface->h };
WGPUExtent3D writeSize{ .width = surface->w, .height = surface->h, .depthOrArrayLayers = 1 };
wgpuQueueWriteTexture(context.queue, &imageCopyTexture, surface->data, 4 * surface->w * surface->h, &textureDataLayout, &writeSize);
wgpuQueueSubmit(context.queue, 0, nullptr);
};
@ -228,20 +211,31 @@ void WgRenderSettings::update(WgContext& context, const Fill* fill, const uint8_
{
// setup fill properties
if ((flags & (RenderUpdateFlag::Gradient)) && fill) {
rasterType = WgRenderRasterType::Gradient;
// setup linear fill properties
if (fill->type() == Type::LinearGradient) {
WgShaderTypeLinearGradient linearGradient((LinearGradient*)fill);
bindGroupLinear.initialize(context.device, context.queue, linearGradient);
if (context.allocateBufferUniform(bufferGroupLinear, &linearGradient, sizeof(linearGradient))) {
context.pipelines->layouts.releaseBindGroup(bindGroupLinear);
bindGroupLinear = context.pipelines->layouts.createBindGroupBuffer1Un(bufferGroupLinear);
}
fillType = WgRenderSettingsType::Linear;
} else if (fill->type() == Type::RadialGradient) {
WgShaderTypeRadialGradient radialGradient((RadialGradient*)fill);
bindGroupRadial.initialize(context.device, context.queue, radialGradient);
if (context.allocateBufferUniform(bufferGroupRadial, &radialGradient, sizeof(radialGradient))) {
context.pipelines->layouts.releaseBindGroup(bindGroupRadial);
bindGroupRadial = context.pipelines->layouts.createBindGroupBuffer1Un(bufferGroupRadial);
}
fillType = WgRenderSettingsType::Radial;
}
skip = false;
} else if ((flags & (RenderUpdateFlag::Color)) && !fill) {
rasterType = WgRenderRasterType::Solid;
WgShaderTypeSolidColor solidColor(color);
bindGroupSolid.initialize(context.device, context.queue, solidColor);
if (context.allocateBufferUniform(bufferGroupSolid, &solidColor, sizeof(solidColor))) {
context.pipelines->layouts.releaseBindGroup(bindGroupSolid);
bindGroupSolid = context.pipelines->layouts.createBindGroupBuffer1Un(bufferGroupSolid);
}
fillType = WgRenderSettingsType::Solid;
skip = (color[3] == 0);
}
@ -250,9 +244,13 @@ void WgRenderSettings::update(WgContext& context, const Fill* fill, const uint8_
void WgRenderSettings::release(WgContext& context)
{
bindGroupSolid.release();
bindGroupLinear.release();
bindGroupRadial.release();
context.pipelines->layouts.releaseBindGroup(bindGroupSolid);
context.pipelines->layouts.releaseBindGroup(bindGroupLinear);
context.pipelines->layouts.releaseBindGroup(bindGroupRadial);
context.releaseBuffer(bufferGroupSolid);
context.releaseBuffer(bufferGroupLinear);
context.releaseBuffer(bufferGroupRadial);
};
//***********************************************************************
@ -261,11 +259,26 @@ void WgRenderSettings::release(WgContext& context)
void WgRenderDataPaint::release(WgContext& context)
{
bindGroupPaint.release();
context.pipelines->layouts.releaseBindGroup(bindGroupPaint);
context.releaseBuffer(bufferModelMat);
context.releaseBuffer(bufferBlendSettings);
clips.clear();
};
void WgRenderDataPaint::update(WgContext& context, const tvg::Matrix& transform, tvg::ColorSpace cs, uint8_t opacity)
{
WgShaderTypeMat4x4f modelMat(transform);
WgShaderTypeBlendSettings blendSettings(cs, opacity);
bool bufferModelMatChanged = context.allocateBufferUniform(bufferModelMat, &modelMat, sizeof(modelMat));
bool bufferBlendSettingsChanged = context.allocateBufferUniform(bufferBlendSettings, &blendSettings, sizeof(blendSettings));
if (bufferModelMatChanged || bufferBlendSettingsChanged) {
context.pipelines->layouts.releaseBindGroup(bindGroupPaint);
bindGroupPaint = context.pipelines->layouts.createBindGroupBuffer2Un(bufferModelMat, bufferBlendSettings);
}
}
void WgRenderDataPaint::updateClips(tvg::Array<tvg::RenderData> &clips) {
this->clips.clear();
for (uint32_t i = 0; i < clips.count; i++)
@ -496,8 +509,6 @@ void WgRenderDataPicture::release(WgContext& context)
{
meshData.release(context);
imageData.release(context);
bindGroupPicture.release();
context.pipelines->layouts.releaseBindGroup(bindGroupPicture);
WgRenderDataPaint::release(context);
}
#endif //_TVG_WG_RENDER_DATA_H_

24
src/renderer/wg_engine/tvgWgRenderData.h Normal file → Executable file
View file

@ -20,6 +20,9 @@
* SOFTWARE.
*/
#ifndef _TVG_WG_RENDER_DATA_H_
#define _TVG_WG_RENDER_DATA_H_
#include "tvgWgPipelines.h"
#include "tvgWgGeometry.h"
@ -70,13 +73,18 @@ struct WgImageData {
};
enum class WgRenderSettingsType { None = 0, Solid = 1, Linear = 2, Radial = 3 };
enum class WgRenderRasterType { Solid = 0, Gradient, Image };
struct WgRenderSettings
{
WgBindGroupSolidColor bindGroupSolid{};
WgBindGroupLinearGradient bindGroupLinear{};
WgBindGroupRadialGradient bindGroupRadial{};
WGPUBuffer bufferGroupSolid{};
WGPUBuffer bufferGroupLinear{};
WGPUBuffer bufferGroupRadial{};
WGPUBindGroup bindGroupSolid{};
WGPUBindGroup bindGroupLinear{};
WGPUBindGroup bindGroupRadial{};
WgRenderSettingsType fillType{};
WgRenderRasterType rasterType{};
bool skip{};
void update(WgContext& context, const Fill* fill, const uint8_t* color, const RenderUpdateFlag flags);
@ -85,7 +93,9 @@ struct WgRenderSettings
struct WgRenderDataPaint
{
WgBindGroupPaint bindGroupPaint{};
WGPUBuffer bufferModelMat{};
WGPUBuffer bufferBlendSettings{};
WGPUBindGroup bindGroupPaint{};
RenderRegion viewport{};
float opacity{};
Array<WgRenderDataPaint*> clips;
@ -94,6 +104,7 @@ struct WgRenderDataPaint
virtual void release(WgContext& context);
virtual Type type() { return Type::Undefined; };
void update(WgContext& context, const tvg::Matrix& transform, tvg::ColorSpace cs, uint8_t opacity);
void updateClips(tvg::Array<tvg::RenderData> &clips);
};
@ -133,11 +144,12 @@ public:
struct WgRenderDataPicture: public WgRenderDataPaint
{
WgBindGroupPicture bindGroupPicture{};
WGPUBindGroup bindGroupPicture{};
WgImageData imageData{};
WgMeshData meshData{};
void update(WgContext& context);
void release(WgContext& context) override;
Type type() override { return Type::Picture; };
};
#endif // _TVG_WG_RENDER_DATA_H_

426
src/renderer/wg_engine/tvgWgRenderTarget.cpp Normal file → Executable file
View file

@ -22,407 +22,69 @@
#include "tvgWgRenderTarget.h"
//*****************************************************************************
// render storage
//*****************************************************************************
void WgRenderStorage::initialize(WgContext& context, uint32_t w, uint32_t h, uint32_t samples, WGPUTextureFormat format)
{
release(context);
// store target storage size
this->samples = samples;
width = w * samples;
height = h * samples;
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
texColor = context.createTexture2d(
WGPUTextureUsage_CopySrc |
WGPUTextureUsage_CopyDst |
WGPUTextureUsage_TextureBinding |
WGPUTextureUsage_StorageBinding |
WGPUTextureUsage_RenderAttachment,
format, width, height, "The target texture color");
texStencil = context.createTexture2d(
WGPUTextureUsage_RenderAttachment,
WGPUTextureFormat_Stencil8,
width, height, "The target texture stencil");
assert(texColor);
assert(texStencil);
texViewColor = context.createTextureView2d(texColor, "The target texture view color");
texViewStencil = context.createTextureView2d(texStencil, "The target texture view stencil");
assert(texViewColor);
assert(texViewStencil);
// initialize bind group for blitting
if (format == WGPUTextureFormat_RGBA8Unorm) {
bindGroupTexStorageRgbaRO.initialize(context.device, context.queue, texViewColor);
bindGroupTexStorageRgbaWO.initialize(context.device, context.queue, texViewColor);
}
if (format == WGPUTextureFormat_BGRA8Unorm)
bindGroupTexStorageBgraWO.initialize(context.device, context.queue, texViewColor);
// initialize window binding groups
WgShaderTypeMat4x4f viewMat(w, h);
mBindGroupCanvas.initialize(context.device, context.queue, viewMat);
mPipelines = context.pipelines;
}
void WgRenderStorage::initialize(WgContext& context, uint32_t width, uint32_t height)
{
this->width = width;
this->height = height;
texture = context.createTexStorage(width, height, WGPUTextureFormat_RGBA8Unorm);
texView = context.createTextureView(texture);
bindGroupRead = context.pipelines->layouts.createBindGroupStrorage1RO(texView);
bindGroupWrite = context.pipelines->layouts.createBindGroupStrorage1WO(texView);
bindGroupTexure = context.pipelines->layouts.createBindGroupTexSampled(context.samplerNearest, texView);
}
void WgRenderStorage::release(WgContext& context)
{
mRenderPassEncoder = nullptr;
mBindGroupCanvas.release();
bindGroupTexStorageBgraWO.release();
bindGroupTexStorageRgbaWO.release();
bindGroupTexStorageRgbaRO.release();
context.releaseTextureView(texViewStencil);
context.releaseTextureView(texViewColor);
context.releaseTexture(texStencil);
context.releaseTexture(texColor);
workgroupsCountX = 0;
workgroupsCountY = 0;
context.pipelines->layouts.releaseBindGroup(bindGroupTexure);
context.pipelines->layouts.releaseBindGroup(bindGroupWrite);
context.pipelines->layouts.releaseBindGroup(bindGroupRead);
context.releaseTextureView(texView);
context.releaseTexture(texture);
height = 0;
width = 0;
}
void WgRenderStorage::renderShape(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType)
{
assert(renderData);
assert(mRenderPassEncoder);
if (renderData->strokeFirst) {
drawStroke(context, renderData, blendType);
drawShape(context, renderData, blendType);
} else {
drawShape(context, renderData, blendType);
drawStroke(context, renderData, blendType);
}
}
void WgRenderStorage::renderPicture(WgContext& context, WgRenderDataPicture* renderData, WgPipelineBlendType blendType)
{
assert(renderData);
assert(mRenderPassEncoder);
uint8_t blend = (uint8_t)blendType;
auto& vp = renderData->viewport;
if ((vp.w <= 0) || (vp.h <= 0)) return;
wgpuRenderPassEncoderSetScissorRect(mRenderPassEncoder, vp.x * samples, vp.y * samples, vp.w * samples, vp.h * samples);
wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 0);
mPipelines->image[blend].use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint, renderData->bindGroupPicture);
renderData->meshData.drawImage(context, mRenderPassEncoder);
}
void WgRenderStorage::renderClipPath(WgContext& context, WgRenderDataPaint* renderData) {
assert(renderData);
assert(mRenderPassEncoder);
if (renderData->type() == Type::Shape)
drawShapeClipPath(context, (WgRenderDataShape*)renderData);
else if (renderData->type() == Type::Picture)
drawPictureClipPath(context, (WgRenderDataPicture*)renderData);
}
void WgRenderStorage::drawShape(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType)
{
assert(renderData);
assert(mRenderPassEncoder);
assert(renderData->meshGroupShapes.meshes.count == renderData->meshGroupShapesBBox.meshes.count);
if (renderData->renderSettingsShape.skip) return;
// apply viewport
auto& vp = renderData->viewport;
if ((vp.w <= 0) || (vp.h <= 0)) return;
wgpuRenderPassEncoderSetScissorRect(mRenderPassEncoder, vp.x * samples, vp.y * samples, vp.w * samples, vp.h * samples);
// setup fill rule
wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 0);
if (renderData->fillRule == FillRule::Winding)
mPipelines->fillShapeWinding.use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint);
else
mPipelines->fillShapeEvenOdd.use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint);
// draw to stencil (first pass)
for (uint32_t i = 0; i < renderData->meshGroupShapes.meshes.count; i++)
renderData->meshGroupShapes.meshes[i]->drawFan(context, mRenderPassEncoder);
// fill shape geometry (second pass)
uint8_t blend = (uint8_t)blendType;
WgRenderSettings& settings = renderData->renderSettingsShape;
if (settings.fillType == WgRenderSettingsType::Solid)
mPipelines->solid[blend].use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint, settings.bindGroupSolid);
else if (settings.fillType == WgRenderSettingsType::Linear)
mPipelines->linear[blend].use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint, settings.bindGroupLinear);
else if (settings.fillType == WgRenderSettingsType::Radial)
mPipelines->radial[blend].use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint, settings.bindGroupRadial);
renderData->meshDataBBox.drawFan(context, mRenderPassEncoder);
}
void WgRenderStorage::drawStroke(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType)
{
assert(renderData);
assert(mRenderPassEncoder);
assert(renderData->meshGroupStrokes.meshes.count == renderData->meshGroupStrokesBBox.meshes.count);
if (renderData->renderSettingsStroke.skip) return;
// apply viewport
auto& vp = renderData->viewport;
if ((vp.w <= 0) || (vp.h <= 0)) return;
wgpuRenderPassEncoderSetScissorRect(mRenderPassEncoder, vp.x * samples, vp.y * samples, vp.w * samples, vp.h * samples);
// draw stroke geometry
uint8_t blend = (uint8_t)blendType;
// draw strokes to stencil (first pass)
for (uint32_t i = 0; i < renderData->meshGroupStrokes.meshes.count; i++) {
// draw to stencil (first pass)
wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 255);
mPipelines->fillStroke.use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint);
renderData->meshGroupStrokes.meshes[i]->draw(context, mRenderPassEncoder);
// fill shape (second pass)
wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 0);
WgRenderSettings& settings = renderData->renderSettingsStroke;
if (settings.fillType == WgRenderSettingsType::Solid)
mPipelines->solid[blend].use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint, settings.bindGroupSolid);
else if (settings.fillType == WgRenderSettingsType::Linear)
mPipelines->linear[blend].use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint, settings.bindGroupLinear);
else if (settings.fillType == WgRenderSettingsType::Radial)
mPipelines->radial[blend].use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint, settings.bindGroupRadial);
renderData->meshGroupStrokesBBox.meshes[i]->drawFan(context, mRenderPassEncoder);
}
}
void WgRenderStorage::drawShapeClipPath(WgContext& context, WgRenderDataShape* renderData) {
assert(renderData);
assert(renderData->type() == Type::Shape);
assert(renderData->meshGroupShapes.meshes.count == renderData->meshGroupShapesBBox.meshes.count);
// draw shape geometry
wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 0);
// setup fill rule
if (renderData->fillRule == FillRule::Winding)
mPipelines->fillShapeWinding.use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint);
else
mPipelines->fillShapeEvenOdd.use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint);
// draw to stencil (first pass)
for (uint32_t i = 0; i < renderData->meshGroupShapes.meshes.count; i++)
renderData->meshGroupShapes.meshes[i]->drawFan(context, mRenderPassEncoder);
// fill shape geometry (second pass)
mPipelines->clipMask.use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint);
renderData->meshDataBBox.drawFan(context, mRenderPassEncoder);
}
void WgRenderStorage::drawPictureClipPath(WgContext& context, WgRenderDataPicture* renderData) {
assert(renderData);
assert(renderData->type() == Type::Picture);
assert(mRenderPassEncoder);
wgpuRenderPassEncoderSetStencilReference(mRenderPassEncoder, 0);
mPipelines->clipMask.use(mRenderPassEncoder, mBindGroupCanvas, renderData->bindGroupPaint);
renderData->meshData.drawImage(context, mRenderPassEncoder);
}
void WgRenderStorage::blend(
WgContext& context,
WGPUCommandEncoder commandEncoder,
WgPipelineBlend* pipeline,
WgRenderStorage* texSrc,
WgRenderStorage* texDst,
WgBindGroupBlendMethod* blendMethod,
WgBindGroupOpacity* opacity)
{
assert(commandEncoder);
assert(texSrc);
assert(texDst);
assert(pipeline);
WgBindGroupTexBlend texBlend;
texBlend.initialize(context.device, context.queue, texSrc->texViewColor, texDst->texViewColor, texViewColor);
WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder);
pipeline->use(computePassEncoder, texBlend, *blendMethod, *opacity);
dispatchWorkgroups(computePassEncoder);
endComputePass(computePassEncoder);
}
void WgRenderStorage::blendMask(
WgContext& context,
WGPUCommandEncoder commandEncoder,
WgPipelineBlendMask* pipeline,
WgRenderStorage* texMsk,
WgRenderStorage* texSrc,
WgRenderStorage* texDst,
WgBindGroupBlendMethod* blendMethod,
WgBindGroupOpacity* opacity)
{
assert(commandEncoder);
assert(texSrc);
assert(texMsk);
WgBindGroupTexBlendMask texBlendMask;
texBlendMask.initialize(context.device, context.queue, texSrc->texViewColor, texMsk->texViewColor, texDst->texViewColor, texViewColor);
WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder);
pipeline->use(computePassEncoder, texBlendMask, *blendMethod, *opacity);
dispatchWorkgroups(computePassEncoder);
endComputePass(computePassEncoder);
texBlendMask.release();
};
void WgRenderStorage::maskCompose(
WgContext& context,
WGPUCommandEncoder commandEncoder,
WgRenderStorage* texMsk0,
WgRenderStorage* texMsk1)
{
assert(commandEncoder);
assert(texMsk0);
WgBindGroupTexMaskCompose maskCompose;
maskCompose.initialize(context.device, context.queue, texMsk0->texViewColor, texMsk1->texViewColor, texViewColor);
WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder);
mPipelines->computeMaskCompose.use(computePassEncoder, maskCompose);
dispatchWorkgroups(computePassEncoder);
endComputePass(computePassEncoder);
maskCompose.release();
}
void WgRenderStorage::compose(
WgContext& context,
WGPUCommandEncoder commandEncoder,
WgRenderStorage* texSrc,
WgRenderStorage* texMsk,
WgRenderStorage* texDst,
WgBindGroupCompositeMethod* composeMethod,
WgBindGroupBlendMethod* blendMethod,
WgBindGroupOpacity* opacity)
{
assert(commandEncoder);
assert(texSrc);
assert(texMsk);
WgBindGroupTexCompose texCompose;
texCompose.initialize(context.device, context.queue, texSrc->texViewColor, texMsk->texViewColor, texDst->texViewColor, texViewColor);
WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder);
mPipelines->computeCompose.use(computePassEncoder, texCompose, *composeMethod, *blendMethod, *opacity);
dispatchWorkgroups(computePassEncoder);
endComputePass(computePassEncoder);
texCompose.release();
}
void WgRenderStorage::antialias(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc)
{
assert(commandEncoder);
assert(targetSrc);
WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder);
mPipelines->computeAntiAliasing.use(computePassEncoder, targetSrc->bindGroupTexStorageRgbaRO, bindGroupTexStorageBgraWO);
dispatchWorkgroups(computePassEncoder);
endComputePass(computePassEncoder);
}
void WgRenderStorage::copy(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc)
{
assert(commandEncoder);
assert(targetSrc);
WGPUComputePassEncoder computePassEncoder = beginComputePass(commandEncoder);
mPipelines->computeCopy.use(computePassEncoder, targetSrc->bindGroupTexStorageRgbaRO, bindGroupTexStorageRgbaWO);
dispatchWorkgroups(computePassEncoder);
endComputePass(computePassEncoder);
}
void WgRenderStorage::dispatchWorkgroups(WGPUComputePassEncoder computePassEncoder)
{
assert(computePassEncoder);
wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, workgroupsCountX, workgroupsCountY, 1);
}
void WgRenderStorage::beginRenderPass(WGPUCommandEncoder commandEncoder, bool clear)
{
assert(commandEncoder);
// render pass depth stencil attachment
WGPURenderPassDepthStencilAttachment depthStencilAttachment{};
depthStencilAttachment.view = texViewStencil;
depthStencilAttachment.depthLoadOp = WGPULoadOp_Undefined;
depthStencilAttachment.depthStoreOp = WGPUStoreOp_Undefined;
depthStencilAttachment.depthClearValue = 1.0f;
depthStencilAttachment.depthReadOnly = false;
depthStencilAttachment.stencilLoadOp = WGPULoadOp_Clear;
depthStencilAttachment.stencilStoreOp = WGPUStoreOp_Discard;
depthStencilAttachment.stencilClearValue = 0;
depthStencilAttachment.stencilReadOnly = false;
// render pass color attachment
WGPURenderPassColorAttachment colorAttachment{};
colorAttachment.view = texViewColor;
#ifdef __EMSCRIPTEN__
colorAttachment.depthSlice = WGPU_DEPTH_SLICE_UNDEFINED;
#endif
colorAttachment.resolveTarget = nullptr;
colorAttachment.loadOp = clear ? WGPULoadOp_Clear : WGPULoadOp_Load;
colorAttachment.clearValue = { 0, 0, 0, 0 };
colorAttachment.storeOp = WGPUStoreOp_Store;
// render pass descriptor
WGPURenderPassDescriptor renderPassDesc{};
renderPassDesc.nextInChain = nullptr;
renderPassDesc.label = "The render pass";
renderPassDesc.colorAttachmentCount = 1;
renderPassDesc.colorAttachments = &colorAttachment;
renderPassDesc.depthStencilAttachment = &depthStencilAttachment;
//renderPassDesc.depthStencilAttachment = nullptr;
renderPassDesc.occlusionQuerySet = nullptr;
renderPassDesc.timestampWrites = nullptr;
// begin render pass
mRenderPassEncoder = wgpuCommandEncoderBeginRenderPass(commandEncoder, &renderPassDesc);
}
void WgRenderStorage::endRenderPass()
{
assert(mRenderPassEncoder);
wgpuRenderPassEncoderEnd(mRenderPassEncoder);
wgpuRenderPassEncoderRelease(mRenderPassEncoder);
mRenderPassEncoder = nullptr;
}
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::endComputePass(WGPUComputePassEncoder computePassEncoder)
{
assert(computePassEncoder);
wgpuComputePassEncoderEnd(computePassEncoder);
}
//*****************************************************************************
// render storage pool
//*****************************************************************************
WgRenderStorage* WgRenderStoragePool::allocate(WgContext& context, uint32_t w, uint32_t h, uint32_t samples)
WgRenderStorage* WgRenderStoragePool::allocate(WgContext& context)
{
WgRenderStorage* renderStorage{};
if (mPool.count > 0) {
renderStorage = mPool.last();
mPool.pop();
} else {
renderStorage = new WgRenderStorage;
renderStorage->initialize(context, w, h, samples);
mList.push(renderStorage);
}
return renderStorage;
WgRenderStorage* renderStorage{};
if (pool.count > 0) {
renderStorage = pool.last();
pool.pop();
} else {
renderStorage = new WgRenderStorage;
renderStorage->initialize(context, width, height);
list.push(renderStorage);
}
return renderStorage;
};
void WgRenderStoragePool::free(WgContext& context, WgRenderStorage* renderStorage) {
mPool.push(renderStorage);
void WgRenderStoragePool::free(WgContext& context, WgRenderStorage* renderStorage)
{
pool.push(renderStorage);
};
void WgRenderStoragePool::initialize(WgContext& context, uint32_t width, uint32_t height)
{
this->width = width;
this->height = height;
}
void WgRenderStoragePool::release(WgContext& context)
{
for (uint32_t i = 0; i < mList.count; i++) {
mList[i]->release(context);
delete mList[i];
}
mList.clear();
mPool.clear();
for (uint32_t i = 0; i < list.count; i++) {
list[i]->release(context);
delete list[i];
}
list.clear();
pool.clear();
height = 0;
width = 0;
};

100
src/renderer/wg_engine/tvgWgRenderTarget.h Normal file → Executable file
View file

@ -23,94 +23,34 @@
#ifndef _TVG_WG_RENDER_TARGET_H_
#define _TVG_WG_RENDER_TARGET_H_
#include "tvgWgRenderData.h"
#include "tvgWgPipelines.h"
#include "tvgRender.h"
class WgRenderStorage {
private:
// texture buffers
WgBindGroupCanvas mBindGroupCanvas;
WGPURenderPassEncoder mRenderPassEncoder{};
WgPipelines* mPipelines{}; // external handle
public:
WGPUTexture texColor{};
WGPUTexture texStencil{};
WGPUTextureView texViewColor{};
WGPUTextureView texViewStencil{};
WgBindGroupTextureStorageRgbaRO bindGroupTexStorageRgbaRO;
WgBindGroupTextureStorageRgbaWO bindGroupTexStorageRgbaWO;
WgBindGroupTextureStorageBgraWO bindGroupTexStorageBgraWO;
uint32_t samples{};
struct WgRenderStorage {
WGPUTexture texture{};
WGPUTextureView texView{};
WGPUBindGroup bindGroupRead{};
WGPUBindGroup bindGroupWrite{};
WGPUBindGroup bindGroupTexure{};
uint32_t width{};
uint32_t height{};
uint32_t workgroupsCountX{};
uint32_t workgroupsCountY{};
public:
void initialize(WgContext& context, uint32_t w, uint32_t h, uint32_t samples = 1, WGPUTextureFormat format = WGPUTextureFormat_RGBA8Unorm);
void initialize(WgContext& context, uint32_t width, uint32_t height);
void release(WgContext& context);
void beginRenderPass(WGPUCommandEncoder commandEncoder, bool clear);
void endRenderPass();
void renderShape(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType);
void renderPicture(WgContext& context, WgRenderDataPicture* renderData, WgPipelineBlendType blendType);
void renderClipPath(WgContext& context, WgRenderDataPaint* renderData);
void blend(
WgContext& context,
WGPUCommandEncoder commandEncoder,
WgPipelineBlend* pipeline,
WgRenderStorage* targetSrc,
WgRenderStorage* targetDst,
WgBindGroupBlendMethod* blendMethod,
WgBindGroupOpacity* opacity);
void blendMask(
WgContext& context,
WGPUCommandEncoder commandEncoder,
WgPipelineBlendMask* pipeline,
WgRenderStorage* texMsk,
WgRenderStorage* texSrc,
WgRenderStorage* texDst,
WgBindGroupBlendMethod* blendMethod,
WgBindGroupOpacity* opacity);
void maskCompose(
WgContext& context,
WGPUCommandEncoder commandEncoder,
WgRenderStorage* texMsk0,
WgRenderStorage* texMsk1);
void compose(
WgContext& context,
WGPUCommandEncoder commandEncoder,
WgRenderStorage* texMsk,
WgRenderStorage* texSrc,
WgRenderStorage* texDst,
WgBindGroupCompositeMethod* composeMethod,
WgBindGroupBlendMethod* blendMethod,
WgBindGroupOpacity* opacity);
void antialias(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc);
void copy(WGPUCommandEncoder commandEncoder, WgRenderStorage* targetSrc);
private:
void drawShape(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType);
void drawStroke(WgContext& context, WgRenderDataShape* renderData, WgPipelineBlendType blendType);
void drawShapeClipPath(WgContext& context, WgRenderDataShape* renderData);
void drawPictureClipPath(WgContext& context, WgRenderDataPicture* renderData);
void dispatchWorkgroups(WGPUComputePassEncoder computePassEncoder);
WGPUComputePassEncoder beginComputePass(WGPUCommandEncoder commandEncoder);
void endComputePass(WGPUComputePassEncoder computePassEncoder);
};
class WgRenderStoragePool {
private:
Array<WgRenderStorage*> mList;
Array<WgRenderStorage*> mPool;
Array<WgRenderStorage*> list;
Array<WgRenderStorage*> pool;
uint32_t width{};
uint32_t height{};
public:
WgRenderStorage* allocate(WgContext& context, uint32_t w, uint32_t h, uint32_t samples = 1);
void free(WgContext& context, WgRenderStorage* renderTarget);
void release(WgContext& context);
WgRenderStorage* allocate(WgContext& context);
void free(WgContext& context, WgRenderStorage* renderTarget);
void initialize(WgContext& context, uint32_t width, uint32_t height);
void release(WgContext& context);
};
#endif
#endif // _TVG_WG_RENDER_TARGET_H_

388
src/renderer/wg_engine/tvgWgRenderer.cpp Normal file → Executable file
View file

@ -21,15 +21,12 @@
*/
#include "tvgWgRenderer.h"
#define WG_SSAA_SAMPLES (2)
/************************************************************************/
/* Internal Class Implementation */
/************************************************************************/
#include <iostream>
WgRenderer::WgRenderer()
{
WgGeometryData::gMath = new WgMath();
WgGeometryData::gMath->initialize();
}
@ -37,49 +34,40 @@ WgRenderer::~WgRenderer()
{
release();
mContext.release();
WgGeometryData::gMath->release();
delete WgGeometryData::gMath;
}
void WgRenderer::initialize()
{
mPipelines.initialize(mContext);
mOpacityPool.initialize(mContext);
mBlendMethodPool.initialize(mContext);
mCompositeMethodPool.initialize(mContext);
WgMeshDataGroup::gMeshDataPool = new WgMeshDataPool();
WgGeometryData::gMath = new WgMath();
WgGeometryData::gMath->initialize();
}
void WgRenderer::release()
{
clearDisposes();
WgGeometryData::gMath->release();
delete WgGeometryData::gMath;
disposeObjects();
mContext.releaseTexture(mTexScreen);
mContext.releaseTextureView(mTexViewScreen);
mContext.pipelines->layouts.releaseBindGroup(mBindGroupScreen);
mStorageRoot.release(mContext);
mRenderStoragePool.release(mContext);
mRenderDataShapePool.release(mContext);
WgMeshDataGroup::gMeshDataPool->release(mContext);
delete WgMeshDataGroup::gMeshDataPool;
mCompositorStack.clear();
mRenderStorageStack.clear();
mRenderStoragePool.release(mContext);
mCompositeMethodPool.release(mContext);
mBlendMethodPool.release(mContext);
mOpacityPool.release(mContext);
mRenderStorageRoot.release(mContext);
mRenderStorageMask.release(mContext);
mRenderStorageCopy.release(mContext);
mRenderStorageScreen.release(mContext);
mRenderStorageInterm.release(mContext);
mPipelines.release();
mPipelines.release(mContext);
}
void WgRenderer::clearDisposes()
void WgRenderer::disposeObjects()
{
if (mDisposed.renderDatas.count == 0) return;
if (mDisposeRenderDatas.count == 0) return;
for (auto p = mDisposed.renderDatas.begin(); p < mDisposed.renderDatas.end(); p++) {
for (auto p = mDisposeRenderDatas.begin(); p < mDisposeRenderDatas.end(); p++) {
auto renderData = (WgRenderDataPaint*)(*p);
if (renderData->type() == Type::Shape) {
mRenderDataShapePool.free(mContext, (WgRenderDataShape*)renderData);
@ -87,12 +75,9 @@ void WgRenderer::clearDisposes()
renderData->release(mContext);
}
}
mDisposed.renderDatas.clear();
mDisposeRenderDatas.clear();
}
/************************************************************************/
/* External Class Implementation */
/************************************************************************/
RenderData WgRenderer::prepare(const RenderShape& rshape, RenderData data, const Matrix& transform, Array<RenderData>& clips, uint8_t opacity, RenderUpdateFlag flags, bool clipper)
{
@ -106,11 +91,9 @@ RenderData WgRenderer::prepare(const RenderShape& rshape, RenderData data, const
renderDataShape->updateMeshes(mContext, rshape, transform);
}
// update paint settings
// update paint settings
if ((!data) || (flags & (RenderUpdateFlag::Transform | RenderUpdateFlag::Blend))) {
WgShaderTypeMat4x4f modelMat(transform);
WgShaderTypeBlendSettings blendSettings(mTargetSurface.cs, opacity);
renderDataShape->bindGroupPaint.initialize(mContext.device, mContext.queue, modelMat, blendSettings);
renderDataShape->update(mContext, transform, mTargetSurface.cs, opacity);
renderDataShape->fillRule = rshape.rule;
}
@ -145,9 +128,7 @@ RenderData WgRenderer::prepare(Surface* surface, const RenderMesh* mesh, RenderD
renderDataPicture->viewport = mViewport;
renderDataPicture->opacity = opacity;
if (flags & (RenderUpdateFlag::Transform | RenderUpdateFlag::Blend)) {
WgShaderTypeMat4x4f modelMat(transform);
WgShaderTypeBlendSettings blendSettings(surface->cs, opacity);
renderDataPicture->bindGroupPaint.initialize(mContext.device, mContext.queue, modelMat, blendSettings);
renderDataPicture->update(mContext, transform, surface->cs, opacity);
}
// update image data
@ -155,13 +136,13 @@ RenderData WgRenderer::prepare(Surface* surface, const RenderMesh* mesh, RenderD
WgGeometryData geometryData;
if (mesh->triangleCnt == 0) geometryData.appendImageBox(surface->w, surface->h);
else geometryData.appendMesh(mesh);
mContext.pipelines->layouts.releaseBindGroup(renderDataPicture->bindGroupPicture);
renderDataPicture->meshData.release(mContext);
renderDataPicture->meshData.update(mContext, &geometryData);
renderDataPicture->imageData.update(mContext, surface);
renderDataPicture->bindGroupPicture.initialize(
mContext.device, mContext.queue,
mContext.samplerLinear,
renderDataPicture->imageData.textureView);
renderDataPicture->bindGroupPicture = mContext.pipelines->layouts.createBindGroupTexSampled(
mContext.samplerLinear, renderDataPicture->imageData.textureView
);
}
// store clips data
@ -172,151 +153,47 @@ RenderData WgRenderer::prepare(Surface* surface, const RenderMesh* mesh, RenderD
bool WgRenderer::preRender()
{
// push rot render storage to the render tree stack
assert(mRenderStorageStack.count == 0);
mRenderStorageStack.push(&mStorageRoot);
// create command encoder for drawing
WGPUCommandEncoderDescriptor commandEncoderDesc{};
commandEncoderDesc.nextInChain = nullptr;
commandEncoderDesc.label = "The command encoder";
mCommandEncoder = wgpuDeviceCreateCommandEncoder(mContext.device, &commandEncoderDesc);
mRenderStorageStack.push(&mRenderStorageRoot);
//mRenderStorageRoot.clear(mCommandEncoder);
mRenderStorageRoot.beginRenderPass(mCommandEncoder, true);
// start root render pass
mCompositor.beginRenderPass(mCommandEncoder, mRenderStorageStack.last(), true);
return true;
}
void WgRenderer::renderClipPath(Array<WgRenderDataPaint*>& clips)
{
for (uint32_t i = 0; i < clips.count; i++) {
renderClipPath(clips[i]->clips);
// render image to render target
mRenderStorageInterm.beginRenderPass(mCommandEncoder, true);
mRenderStorageInterm.renderClipPath(mContext, clips[i]);
mRenderStorageInterm.endRenderPass();
mRenderStorageCopy.copy(mCommandEncoder, &mRenderStorageMask);
mRenderStorageMask.maskCompose(mContext, mCommandEncoder, &mRenderStorageInterm, &mRenderStorageCopy);
}
}
bool WgRenderer::renderShape(RenderData data)
{
// get current render storage
WgRenderDataShape *dataShape = (WgRenderDataShape*)data;
if (dataShape->opacity == 0) return 0;
WgPipelineBlendType blendType = WgPipelines::blendMethodToBlendType(mBlendMethod);
WgRenderStorage* renderStorage = mRenderStorageStack.last();
assert(renderStorage);
// use masked blend
if (dataShape->clips.count > 0) {
// terminate current render pass
renderStorage->endRenderPass();
// render clip path
mRenderStorageMask.beginRenderPass(mCommandEncoder, true);
mRenderStorageMask.renderClipPath(mContext, dataShape->clips[0]);
mRenderStorageMask.endRenderPass();
renderClipPath(dataShape->clips);
// render image to render target
mRenderStorageInterm.beginRenderPass(mCommandEncoder, true);
mRenderStorageInterm.renderShape(mContext, dataShape, blendType);
mRenderStorageInterm.endRenderPass();
// blend shape with current render storage
WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod);
WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, 255);
WgPipelineBlendMask* pipeline = &mContext.pipelines->computeBlendSolidMask;
if (dataShape->renderSettingsShape.fillType != WgRenderSettingsType::Solid)
pipeline = &mContext.pipelines->computeBlendGradientMask;
mRenderStorageCopy.copy(mCommandEncoder, renderStorage);
renderStorage->blendMask(mContext, mCommandEncoder,
pipeline, &mRenderStorageMask, &mRenderStorageInterm, &mRenderStorageCopy, blendMethod, opacity);
// restore current render pass
renderStorage->beginRenderPass(mCommandEncoder, false);
// use hardware blend
} else if (WgPipelines::isBlendMethodSupportsHW(mBlendMethod)) {
renderStorage->renderShape(mContext, dataShape, blendType);
// use custom blend
} else {
// terminate current render pass
renderStorage->endRenderPass();
// render image to render target
mRenderStorageInterm.beginRenderPass(mCommandEncoder, true);
mRenderStorageInterm.renderShape(mContext, dataShape, blendType);
mRenderStorageInterm.endRenderPass();
// blend shape with current render storage
WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod);
WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, 255);
WgPipelineBlend* pipeline = &mContext.pipelines->computeBlendSolid;
if (dataShape->renderSettingsShape.fillType != WgRenderSettingsType::Solid)
pipeline = &mContext.pipelines->computeBlendGradient;
mRenderStorageCopy.copy(mCommandEncoder, renderStorage);
renderStorage->blend(mContext, mCommandEncoder,
pipeline, &mRenderStorageInterm, &mRenderStorageCopy, blendMethod, opacity);
// restore current render pass
renderStorage->beginRenderPass(mCommandEncoder, false);
}
// temporary simple render data to the current render target
mCompositor.renderShape(mContext, (WgRenderDataShape*)data, mBlendMethod);
return true;
}
bool WgRenderer::renderImage(RenderData data)
{
// get current render storage
WgRenderDataPicture *dataPicture = (WgRenderDataPicture*)data;
// get current render storage
WgPipelineBlendType blendType = WgPipelines::blendMethodToBlendType(mBlendMethod);
WgRenderStorage* renderStorage = mRenderStorageStack.last();
assert(renderStorage);
// use masked blend
if (dataPicture->clips.count > 0) {
// terminate current render pass
renderStorage->endRenderPass();
// render clip path
mRenderStorageMask.beginRenderPass(mCommandEncoder, true);
mRenderStorageMask.renderClipPath(mContext, dataPicture->clips[0]);
mRenderStorageMask.endRenderPass();
renderClipPath(dataPicture->clips);
// render image to render target
mRenderStorageInterm.beginRenderPass(mCommandEncoder, true);
mRenderStorageInterm.renderPicture(mContext, dataPicture, blendType);
mRenderStorageInterm.endRenderPass();
// blend shape with current render storage
WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod);
WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, 255);
WgPipelineBlendMask* pipeline = &mContext.pipelines->computeBlendImageMask;
mRenderStorageCopy.copy(mCommandEncoder, renderStorage);
renderStorage->blendMask(mContext, mCommandEncoder,
pipeline, &mRenderStorageMask, &mRenderStorageInterm, &mRenderStorageCopy, blendMethod, opacity);
// restore current render pass
renderStorage->beginRenderPass(mCommandEncoder, false);
// use hardware blend
} else if (WgPipelines::isBlendMethodSupportsHW(mBlendMethod))
renderStorage->renderPicture(mContext, dataPicture, blendType);
// use custom blend
else {
// terminate current render pass
renderStorage->endRenderPass();
// render image to render target
mRenderStorageInterm.beginRenderPass(mCommandEncoder, true);
mRenderStorageInterm.renderPicture(mContext, dataPicture, blendType);
mRenderStorageInterm.endRenderPass();
// blend shape with current render storage
WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod);
WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, 255);
WgPipelineBlend* pipeline = &mContext.pipelines->computeBlendImage;
mRenderStorageCopy.copy(mCommandEncoder, renderStorage);
renderStorage->blend(mContext, mCommandEncoder,
pipeline, &mRenderStorageInterm, &mRenderStorageCopy, blendMethod, opacity);
// restore current render pass
renderStorage->beginRenderPass(mCommandEncoder, false);
}
// temporary simple render data to the current render target
mCompositor.renderImage(mContext, (WgRenderDataPicture*)data, mBlendMethod);
return true;
}
bool WgRenderer::postRender()
{
mRenderStorageRoot.endRenderPass();
mRenderStorageStack.pop();
mContext.executeCommandEncoder(mCommandEncoder);
// end root render pass
mCompositor.endRenderPass();
// release command encoder
const WGPUCommandBufferDescriptor commandBufferDesc{};
WGPUCommandBuffer commandsBuffer = wgpuCommandEncoderFinish(mCommandEncoder, &commandBufferDesc);
wgpuQueueSubmit(mContext.queue, 1, &commandsBuffer);
wgpuCommandBufferRelease(commandsBuffer);
wgpuCommandEncoderRelease(mCommandEncoder);
// pop root render storage to the render tree stack
mRenderStorageStack.pop();
assert(mRenderStorageStack.count == 0);
return true;
}
@ -324,8 +201,8 @@ bool WgRenderer::postRender()
void WgRenderer::dispose(RenderData data) {
auto renderData = (WgRenderDataPaint*)data;
if (renderData) {
ScopedLock lock(mDisposed.key);
mDisposed.renderDatas.push(data);
ScopedLock lock(mDisposeKey);
mDisposeRenderDatas.push(data);
}
}
@ -375,28 +252,35 @@ bool WgRenderer::clear()
bool WgRenderer::sync()
{
clearDisposes();
WGPUSurfaceTexture backBuffer{};
wgpuSurfaceGetCurrentTexture(mContext.surface, &backBuffer);
disposeObjects();
// get current texture
WGPUSurfaceTexture surfaceTexture{};
wgpuSurfaceGetCurrentTexture(mContext.surface, &surfaceTexture);
WGPUCommandEncoderDescriptor commandEncoderDesc{};
commandEncoderDesc.nextInChain = nullptr;
commandEncoderDesc.label = "The command encoder";
// create command encoder
const WGPUCommandEncoderDescriptor commandEncoderDesc{};
WGPUCommandEncoder commandEncoder = wgpuDeviceCreateCommandEncoder(mContext.device, &commandEncoderDesc);
mRenderStorageScreen.antialias(commandEncoder, &mRenderStorageRoot);
// copy render to screen with conversion rgba to bgra (screen format)
const WGPUComputePassDescriptor computePassDescriptor{};
WGPUComputePassEncoder computePassEncoder = wgpuCommandEncoderBeginComputePass(commandEncoder, &computePassDescriptor);
wgpuComputePassEncoderSetBindGroup(computePassEncoder, 0, mStorageRoot.bindGroupRead, 0, nullptr);
wgpuComputePassEncoderSetBindGroup(computePassEncoder, 1, mBindGroupScreen, 0, nullptr);
wgpuComputePassEncoderSetPipeline(computePassEncoder, mContext.pipelines->copy);
wgpuComputePassEncoderDispatchWorkgroups(computePassEncoder, (mStorageRoot.width + 7) / 8, (mStorageRoot.height + 7) / 8, 1);
wgpuComputePassEncoderEnd(computePassEncoder);
WGPUImageCopyTexture source{};
source.texture = mRenderStorageScreen.texColor;
WGPUImageCopyTexture dest{};
dest.texture = backBuffer.texture;
WGPUExtent3D copySize{};
copySize.width = mTargetSurface.w;
copySize.height = mTargetSurface.h;
copySize.depthOrArrayLayers = 1;
wgpuCommandEncoderCopyTextureToTexture(commandEncoder, &source, &dest, &copySize);
// copy dst storage to temporary read only storage
const WGPUImageCopyTexture texSrc { .texture = mTexScreen };
const WGPUImageCopyTexture texDst { .texture = surfaceTexture.texture };
const WGPUExtent3D copySize { .width = mTargetSurface.w, .height = mTargetSurface.h, .depthOrArrayLayers = 1 };
wgpuCommandEncoderCopyTextureToTexture(commandEncoder, &texSrc, &texDst, &copySize);
mContext.executeCommandEncoder(commandEncoder);
// release command encoder
const WGPUCommandBufferDescriptor commandBufferDesc{};
WGPUCommandBuffer commandsBuffer = wgpuCommandEncoderFinish(commandEncoder, &commandBufferDesc);
wgpuQueueSubmit(mContext.queue, 1, &commandsBuffer);
wgpuCommandBufferRelease(commandsBuffer);
wgpuCommandEncoderRelease(commandEncoder);
return true;
@ -413,114 +297,88 @@ bool WgRenderer::target(WGPUInstance instance, WGPUSurface surface, uint32_t w,
mContext.initialize(instance, surface);
// configure surface
WGPUSurfaceConfiguration surfaceConfiguration{};
surfaceConfiguration.nextInChain = nullptr;
surfaceConfiguration.device = mContext.device;
surfaceConfiguration.format = WGPUTextureFormat_BGRA8Unorm;
surfaceConfiguration.usage = WGPUTextureUsage_CopyDst;
surfaceConfiguration.viewFormatCount = 0;
surfaceConfiguration.viewFormats = nullptr;
surfaceConfiguration.alphaMode = WGPUCompositeAlphaMode_Auto;
surfaceConfiguration.width = w;
surfaceConfiguration.height = h;
#ifdef __EMSCRIPTEN__
surfaceConfiguration.presentMode = WGPUPresentMode_Fifo;
#else
surfaceConfiguration.presentMode = WGPUPresentMode_Immediate;
#endif
wgpuSurfaceConfigure(mContext.surface, &surfaceConfiguration);
WGPUSurfaceConfiguration surfaceConfiguration {
.device = mContext.device,
.format = WGPUTextureFormat_BGRA8Unorm,
.usage = WGPUTextureUsage_CopyDst,
.width = w, .height = h,
#ifdef __EMSCRIPTEN__
.presentMode = WGPUPresentMode_Fifo,
#else
.presentMode = WGPUPresentMode_Immediate
#endif
};
wgpuSurfaceConfigure(surface, &surfaceConfiguration);
initialize();
mRenderStorageInterm.initialize(mContext, w, h, WG_SSAA_SAMPLES);
mRenderStorageMask.initialize(mContext, w, h, WG_SSAA_SAMPLES);
mRenderStorageCopy.initialize(mContext, w, h, WG_SSAA_SAMPLES);
mRenderStorageRoot.initialize(mContext, w, h, WG_SSAA_SAMPLES);
mRenderStorageScreen.initialize(mContext, w, h, 1, WGPUTextureFormat_BGRA8Unorm);
mRenderStoragePool.initialize(mContext, w, h);
mStorageRoot.initialize(mContext, w, h);
mCompositor.initialize(mContext, w, h);
// screen buffer
mTexScreen = mContext.createTexStorage(w, h, WGPUTextureFormat_BGRA8Unorm);
mTexViewScreen = mContext.createTextureView(mTexScreen);
mBindGroupScreen = mContext.pipelines->layouts.createBindGroupScreen1WO(mTexViewScreen);
return true;
}
Compositor* WgRenderer::target(TVG_UNUSED const RenderRegion& region, TVG_UNUSED ColorSpace cs)
{
mCompositorStack.push(new WgCompositor);
mCompositorStack.push(new WgCompose);
return mCompositorStack.last();
}
bool WgRenderer::beginComposite(TVG_UNUSED Compositor* cmp, TVG_UNUSED CompositeMethod method, TVG_UNUSED uint8_t opacity)
bool WgRenderer::beginComposite(Compositor* cmp, CompositeMethod method, uint8_t opacity)
{
// save current composition settings
WgCompositor *comp = (WgCompositor*)cmp;
comp->method = method;
comp->opacity = opacity;
comp->blendMethod = mBlendMethod;
WgCompose* compose = (WgCompose *)cmp;
compose->method = method;
compose->opacity = opacity;
compose->blend = mBlendMethod;
// end current render pass
mRenderStorageStack.last()->endRenderPass();
// allocate new render storage and push it to top of render tree
WgRenderStorage* renderStorage = mRenderStoragePool.allocate(mContext, mTargetSurface.w, mTargetSurface.h, WG_SSAA_SAMPLES);
mRenderStorageStack.push(renderStorage);
// begin last render pass
mRenderStorageStack.last()->beginRenderPass(mCommandEncoder, true);
mCompositor.endRenderPass();
// allocate new render storage and push to the render tree stack
WgRenderStorage* storage = mRenderStoragePool.allocate(mContext);
mRenderStorageStack.push(storage);
// begin newly added render pass
mCompositor.beginRenderPass(mCommandEncoder, mRenderStorageStack.last(), true);
return true;
}
bool WgRenderer::endComposite(TVG_UNUSED Compositor* cmp)
bool WgRenderer::endComposite(Compositor* cmp)
{
WgCompositor *comp = (WgCompositor*)cmp;
// get current composition settings
WgCompose* comp = (WgCompose *)cmp;
// end current render pass
mCompositor.endRenderPass();
// finish scene blending
if (comp->method == CompositeMethod::None) {
// end current render pass
mRenderStorageStack.last()->endRenderPass();
// get two last render targets
WgRenderStorage* renderStorageSrc = mRenderStorageStack.last();
// get source and destination render storages
WgRenderStorage* src = mRenderStorageStack.last();
mRenderStorageStack.pop();
// blent scene to current render storage
WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, comp->blendMethod);
WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, comp->opacity);
mRenderStorageCopy.copy(mCommandEncoder, mRenderStorageStack.last());
mRenderStorageStack.last()->blend(mContext, mCommandEncoder,
&mContext.pipelines->computeBlendImage,
renderStorageSrc, &mRenderStorageCopy, blendMethod, opacity);
WgRenderStorage* dst = mRenderStorageStack.last();
// apply blend
mCompositor.blend(mCommandEncoder, src, dst, comp->opacity, comp->blend, WgRenderRasterType::Image);
// back render targets to the pool
mRenderStoragePool.free(mContext, renderStorageSrc);
// begin last render pass
mRenderStorageStack.last()->beginRenderPass(mCommandEncoder, false);
} else {
// end current render pass
mRenderStorageStack.last()->endRenderPass();
mRenderStoragePool.free(mContext, src);
} else { // finish composition
// get source, mask and destination render storages
WgRenderStorage* renderStorageSrc = mRenderStorageStack.last();
WgRenderStorage* src = mRenderStorageStack.last();
mRenderStorageStack.pop();
WgRenderStorage* renderStorageMsk = mRenderStorageStack.last();
WgRenderStorage* msk = mRenderStorageStack.last();
mRenderStorageStack.pop();
WgRenderStorage* renderStorageDst = mRenderStorageStack.last();
// get compose, blend and opacity settings
WgBindGroupCompositeMethod* composeMethod = mCompositeMethodPool.allocate(mContext, comp->method);
WgBindGroupBlendMethod* blendMethod = mBlendMethodPool.allocate(mContext, mBlendMethod);
WgBindGroupOpacity* opacity = mOpacityPool.allocate(mContext, comp->opacity);
// compose and blend
// dest = blend(dest, compose(src, msk, composeMethod), blendMethod, opacity)
mRenderStorageCopy.copy(mCommandEncoder, renderStorageDst);
renderStorageDst->compose(mContext, mCommandEncoder,
renderStorageSrc, renderStorageMsk, &mRenderStorageCopy,
composeMethod, blendMethod, opacity);
WgRenderStorage* dst = mRenderStorageStack.last();
// apply composition
mCompositor.compose(mCommandEncoder, src, msk, dst,comp->method);
// back render targets to the pool
mRenderStoragePool.free(mContext, renderStorageSrc);
mRenderStoragePool.free(mContext, renderStorageMsk);
// begin last render pass
mRenderStorageStack.last()->beginRenderPass(mCommandEncoder, false);
mRenderStoragePool.free(mContext, src);
mRenderStoragePool.free(mContext, msk);
}
// begin previous render pass
mCompositor.beginRenderPass(mCommandEncoder, mRenderStorageStack.last(), false);
// delete current compositor
// delete current compositor settings
delete mCompositorStack.last();
mCompositorStack.pop();

37
src/renderer/wg_engine/tvgWgRenderer.h Normal file → Executable file
View file

@ -23,7 +23,7 @@
#ifndef _TVG_WG_RENDERER_H_
#define _TVG_WG_RENDERER_H_
#include "tvgWgRenderTarget.h"
#include "tvgWgCompositor.h"
class WgRenderer : public RenderMethod
{
@ -61,37 +61,32 @@ private:
~WgRenderer();
void initialize();
void release();
void clearDisposes();
void renderClipPath(Array<WgRenderDataPaint*>& clips);
void disposeObjects();
WGPUCommandEncoder mCommandEncoder{};
WgRenderStorage mRenderStorageInterm; // intermidiate buffer to render
WgRenderStorage mRenderStorageCopy; // copy of destination target (blend and compostition)
WgRenderStorage mRenderStorageMask; // buffer to render mask
WgRenderStorage mRenderStorageRoot; // root render storage
WgRenderStorage mRenderStorageScreen; // storage with data after antializing
WgRenderStoragePool mRenderStoragePool; // pool to hold render tree storages
WgBindGroupOpacityPool mOpacityPool; // opacity, blend methods and composite methods pool
WgBindGroupBlendMethodPool mBlendMethodPool;
WgBindGroupCompositeMethodPool mCompositeMethodPool;
WgRenderDataShapePool mRenderDataShapePool; // render data shpes pool
WgRenderDataShapePool mRenderDataShapePool;
// render tree stacks
Array<Compositor*> mCompositorStack;
WgRenderStorage mStorageRoot;
Array<WgCompose*> mCompositorStack;
Array<WgRenderStorage*> mRenderStorageStack;
WgRenderStoragePool mRenderStoragePool;
WgContext mContext;
WgPipelines mPipelines;
WgCompositor mCompositor;
// screen buffer
WGPUTexture mTexScreen{};
WGPUTextureView mTexViewScreen{};
WGPUBindGroup mBindGroupScreen{};
Surface mTargetSurface;
BlendMethod mBlendMethod{};
RenderRegion mViewport{};
BlendMethod mBlendMethod{}; // current blend method
// disposed resources, they should be released on synced call.
struct {
Array<RenderData> renderDatas{};
Key key;
} mDisposed;
Array<RenderData> mDisposeRenderDatas{};
Key mDisposeKey{};
};
#endif /* _TVG_WG_RENDERER_H_ */

583
src/renderer/wg_engine/tvgWgShaderSrc.cpp Normal file → Executable file
View file

@ -26,10 +26,10 @@
#define WG_SHADER_SOURCE(...) #__VA_ARGS__
//************************************************************************
// shader pipeline fill
// graphics shader source: stencil
//************************************************************************
const char* cShaderSource_PipelineFill = WG_SHADER_SOURCE(
const char* cShaderSrc_Stencil = WG_SHADER_SOURCE(
// vertex input
struct VertexInput {
@location(0) position: vec2f
@ -41,8 +41,8 @@ struct VertexOutput {
};
// uniforms
@group(0) @binding(0) var<uniform> uViewMat : mat4x4f;
@group(1) @binding(0) var<uniform> uModelMat : mat4x4f;
@group(0) @binding(0) var<uniform> uViewMat : mat4x4f;
@group(1) @binding(0) var<uniform> uModelMat : mat4x4f;
@vertex
fn vs_main(in: VertexInput) -> VertexOutput {
@ -59,10 +59,10 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f {
);
//************************************************************************
// shader pipeline solid
// graphics shader source: solid
//************************************************************************
const char* cShaderSource_PipelineSolid = WG_SHADER_SOURCE(
const char* cShaderSrc_Solid = WG_SHADER_SOURCE(
// vertex input
struct VertexInput {
@location(0) position: vec2f
@ -109,10 +109,10 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f {
);
//************************************************************************
// shader pipeline linear
// graphics shader source: linear
//************************************************************************
const char* cShaderSource_PipelineLinear = WG_SHADER_SOURCE(
const char* cShaderSrc_Linear = WG_SHADER_SOURCE(
// vertex input
struct VertexInput {
@location(0) position: vec2f
@ -149,7 +149,7 @@ struct VertexOutput {
// uniforms
@group(0) @binding(0) var<uniform> uViewMat : mat4x4f;
@group(1) @binding(0) var<uniform> uModelMat : mat4x4f;
@group(1) @binding(1) var<uniform> uBlendSettings : BlendSettings;
@group(1) @binding(1) var<uniform> uBlendSettings : BlendSettings;
@group(2) @binding(0) var<uniform> uLinearGradient : LinearGradient;
@vertex
@ -212,10 +212,10 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f {
);
//************************************************************************
// shader pipeline radial
// graphics shader source: radial
//************************************************************************
const char* cShaderSource_PipelineRadial = WG_SHADER_SOURCE(
const char* cShaderSrc_Radial = WG_SHADER_SOURCE(
// vertex input
struct VertexInput {
@location(0) position: vec2f
@ -309,10 +309,10 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f {
);
//************************************************************************
// cShaderSource_PipelineImage
// graphics shader source: image
//************************************************************************
const char* cShaderSource_PipelineImage = WG_SHADER_SOURCE(
const char* cShaderSrc_Image = WG_SHADER_SOURCE(
// vertex input
struct VertexInput {
@location(0) position: vec2f,
@ -365,296 +365,329 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f {
);
//************************************************************************
// cShaderSource_PipelineComputeBlend
// compute shaders blend
//************************************************************************
const std::string strBlendShaderHeader = WG_SHADER_SOURCE(
@group(0) @binding(0) var imageSrc : texture_storage_2d<rgba8unorm, read>;
@group(0) @binding(1) var imageDst : texture_storage_2d<rgba8unorm, read>;
@group(0) @binding(2) var imageTrg : texture_storage_2d<rgba8unorm, write>;
@group(1) @binding(0) var<uniform> blendMethod : u32;
@group(2) @binding(0) var<uniform> opacity : f32;
@compute @workgroup_size(8, 8)
fn cs_main( @builtin(global_invocation_id) id: vec3u) {
let texSize = textureDimensions(imageSrc);
if ((id.x >= texSize.x) || (id.y >= texSize.y)) { return; };
let colorSrc= textureLoad(imageSrc, id.xy);
if (colorSrc.a == 0.0) { return; }
let colorDst = textureLoad(imageDst, id.xy);
var One: vec3f = vec3(1.0);
var So: f32 = opacity;
var Sc: vec3f = colorSrc.rgb;
var Sa: f32 = colorSrc.a;
var Dc: vec3f = colorDst.rgb;
var Da: f32 = colorDst.a;
var Rc: vec3f = colorDst.rgb;
var Ra: f32 = 1.0;
);
const std::string strBlendMaskShaderHeader = WG_SHADER_SOURCE(
@group(0) @binding(0) var imageSrc : texture_storage_2d<rgba8unorm, read>;
@group(0) @binding(1) var imageMsk : texture_storage_2d<rgba8unorm, read>;
@group(0) @binding(2) var imageDst : texture_storage_2d<rgba8unorm, read>;
@group(0) @binding(3) var imageTrg : texture_storage_2d<rgba8unorm, write>;
@group(1) @binding(0) var<uniform> blendMethod : u32;
@group(2) @binding(0) var<uniform> opacity : f32;
@compute @workgroup_size(8, 8)
fn cs_main( @builtin(global_invocation_id) id: vec3u) {
let texSize = textureDimensions(imageSrc);
if ((id.x >= texSize.x) || (id.y >= texSize.y)) { return; };
let colorMsk = textureLoad(imageMsk, id.xy);
if (colorMsk.a == 0.0) { return; }
let colorSrc = textureLoad(imageSrc, id.xy);
if (colorSrc.a == 0.0) { return; }
let colorDst = textureLoad(imageDst, id.xy);
var One: vec3f = vec3(1.0);
var So: f32 = opacity;
var Sc: vec3f = colorSrc.rgb;
var Sa: f32 = colorSrc.a;
var Dc: vec3f = colorDst.rgb;
var Da: f32 = colorDst.a;
var Rc: vec3f = colorDst.rgb;
var Ra: f32 = 1.0;
);
const std::string strBlendShaderPreConditionsGradient = WG_SHADER_SOURCE(
Sc = Sc + Dc.rgb * (1.0 - Sa);
Sa = Sa + Da * (1.0 - Sa);
);
const std::string strBlendShaderPreConditionsImage = WG_SHADER_SOURCE(
Sc = Sc * So;
Sa = Sa * So;
);
const std::string strBlendShaderBlendMethod = WG_SHADER_SOURCE(
switch blendMethod {
/* Normal */ case 0u: {
Rc = Sc + Dc * (1.0 - Sa);
Ra = Sa + Da * (1.0 - Sa);
}
/* Add */ case 1u: { Rc = Sc + Dc; }
/* Screen */ case 2u: { Rc = Sc + Dc - Sc * Dc; }
/* Multiply */ case 3u: { Rc = Sc * Dc; }
/* Overlay */ case 4u: {
Rc.r = select(1.0 - min(1.0, 2 * (1 - Sc.r) * (1 - Dc.r)), min(1.0, 2 * Sc.r * Dc.r), (Dc.r < 0.5));
Rc.g = select(1.0 - min(1.0, 2 * (1 - Sc.g) * (1 - Dc.g)), min(1.0, 2 * Sc.g * Dc.g), (Dc.g < 0.5));
Rc.b = select(1.0 - min(1.0, 2 * (1 - Sc.b) * (1 - Dc.b)), min(1.0, 2 * Sc.b * Dc.b), (Dc.b < 0.5));
}
/* Difference */ case 5u: { Rc = abs(Dc - Sc); }
/* Exclusion */ case 6u: { Rc = min(One, Sc + Dc - min(One, 2 * Sc * Dc)); }
/* SrcOver */ case 7u: { Rc = Sc; Ra = Sa; }
/* Darken */ case 8u: { Rc = min(Sc, Dc); }
/* Lighten */ case 9u: { Rc = max(Sc, Dc); }
/* ColorDodge */ case 10u: {
Rc.r = select(Dc.r, (Dc.r * 255.0 / (255.0 - Sc.r * 255.0))/255.0, (1.0 - Sc.r > 0.0));
Rc.g = select(Dc.g, (Dc.g * 255.0 / (255.0 - Sc.g * 255.0))/255.0, (1.0 - Sc.g > 0.0));
Rc.b = select(Dc.b, (Dc.b * 255.0 / (255.0 - Sc.b * 255.0))/255.0, (1.0 - Sc.b > 0.0));
}
/* ColorBurn */ case 11u: {
Rc.r = select(1.0 - Dc.r, (255.0 - (255.0 - Dc.r * 255.0) / (Sc.r * 255.0)) / 255.0, (Sc.r > 0.0));
Rc.g = select(1.0 - Dc.g, (255.0 - (255.0 - Dc.g * 255.0) / (Sc.g * 255.0)) / 255.0, (Sc.g > 0.0));
Rc.b = select(1.0 - Dc.b, (255.0 - (255.0 - Dc.b * 255.0) / (Sc.b * 255.0)) / 255.0, (Sc.b > 0.0));
}
/* HardLight */ case 12u: {
Rc.r = select(1.0 - min(1.0, 2 * (1 - Sc.r) * (1 - Dc.r)), min(1.0, 2 * Sc.r * Dc.r), (Sc.r < 0.5));
Rc.g = select(1.0 - min(1.0, 2 * (1 - Sc.g) * (1 - Dc.g)), min(1.0, 2 * Sc.g * Dc.g), (Sc.g < 0.5));
Rc.b = select(1.0 - min(1.0, 2 * (1 - Sc.b) * (1 - Dc.b)), min(1.0, 2 * Sc.b * Dc.b), (Sc.b < 0.5));
}
/* SoftLight */ case 13u: { Rc = min(One, (One - 2 * Sc) * Dc * Dc + 2.0 * Sc * Dc); }
default: {
Rc = Sc + Dc * (1.0 - Sa);
Ra = Sa + Da * (1.0 - Sa);
}
}
);
const std::string strBlendShaderPostConditionsGradient = WG_SHADER_SOURCE(
// nothing
);
const std::string strBlendShaderPostConditionsImage = WG_SHADER_SOURCE(
Rc = select(mix(Dc, Rc, Sa), Rc, blendMethod == 0);
Ra = select(mix(Da, Ra, Sa), Ra, blendMethod == 0);
);
const std::string strBlendShaderFooter = WG_SHADER_SOURCE(
textureStore(imageTrg, id.xy, vec4(Rc, Ra));
}
);
// pipeline shader modules blend solid
const std::string strComputeBlendSolid =
strBlendShaderHeader +
strBlendShaderBlendMethod +
strBlendShaderFooter;
const char* cShaderSource_PipelineComputeBlendSolid = strComputeBlendSolid.c_str();
// pipeline shader modules blend gradient
const std::string strComputeBlendGradient =
strBlendShaderHeader +
strBlendShaderPreConditionsGradient +
strBlendShaderBlendMethod +
strBlendShaderPostConditionsGradient +
strBlendShaderFooter;
const char* cShaderSource_PipelineComputeBlendGradient = strComputeBlendGradient.c_str();
// pipeline shader modules blend image
const std::string strComputeBlendImage =
strBlendShaderHeader +
strBlendShaderPreConditionsImage +
strBlendShaderBlendMethod +
strBlendShaderPostConditionsImage +
strBlendShaderFooter;
const char* cShaderSource_PipelineComputeBlendImage = strComputeBlendImage.c_str();
// pipeline shader modules blend solid mask
const std::string strComputeBlendSolidMask =
strBlendMaskShaderHeader +
strBlendShaderBlendMethod +
strBlendShaderFooter;
const char* cShaderSource_PipelineComputeBlendSolidMask = strComputeBlendSolidMask.c_str();
// pipeline shader modules blend gradient mask
const std::string strComputeBlendGradientMask =
strBlendMaskShaderHeader +
strBlendShaderPreConditionsGradient +
strBlendShaderBlendMethod +
strBlendShaderPostConditionsGradient +
strBlendShaderFooter;
const char* cShaderSource_PipelineComputeBlendGradientMask = strComputeBlendGradientMask.c_str();
// pipeline shader modules blend image mask
const std::string strComputeBlendImageMask =
strBlendMaskShaderHeader +
strBlendShaderPreConditionsImage +
strBlendShaderBlendMethod +
strBlendShaderPostConditionsImage +
strBlendShaderFooter;
const char* cShaderSource_PipelineComputeBlendImageMask = strComputeBlendImageMask.c_str();
// pipeline shader modules clear
const char* cShaderSource_PipelineComputeClear = WG_SHADER_SOURCE(
@group(0) @binding(0) var imageDst : texture_storage_2d<rgba8unorm, write>;
@compute @workgroup_size(8, 8)
fn cs_main( @builtin(global_invocation_id) id: vec3u) {
textureStore(imageDst, id.xy, vec4f(0.0, 0.0, 0.0, 0.0));
}
);
// pipeline shader modules copy
const char* cShaderSource_PipelineComputeCopy = WG_SHADER_SOURCE(
@group(0) @binding(0) var imageSrc : texture_storage_2d<rgba8unorm, read>;
@group(1) @binding(0) var imageDst : texture_storage_2d<rgba8unorm, write>;
@compute @workgroup_size(8, 8)
fn cs_main( @builtin(global_invocation_id) id: vec3u) {
textureStore(imageDst, id.xy, textureLoad(imageSrc, id.xy));
}
);
// pipeline shader modules compose
const char* cShaderSource_PipelineComputeMaskCompose = WG_SHADER_SOURCE(
const char* cShaderSrc_MergeMasks = WG_SHADER_SOURCE(
@group(0) @binding(0) var imageMsk0 : texture_storage_2d<rgba8unorm, read>;
@group(0) @binding(1) var imageMsk1 : texture_storage_2d<rgba8unorm, read>;
@group(0) @binding(2) var imageTrg : texture_storage_2d<rgba8unorm, write>;
@group(1) @binding(0) var imageMsk1 : texture_storage_2d<rgba8unorm, read>;
@group(2) @binding(0) var imageTrg : texture_storage_2d<rgba8unorm, write>;
@compute @workgroup_size(8, 8)
fn cs_main( @builtin(global_invocation_id) id: vec3u) {
fn cs_main(@builtin(global_invocation_id) id: vec3u) {
let colorMsk0 = textureLoad(imageMsk0, id.xy);
let colorMsk1 = textureLoad(imageMsk1, id.xy);
textureStore(imageTrg, id.xy, colorMsk0 * colorMsk1);
}
);
// pipeline shader modules compose
const char* cShaderSource_PipelineComputeCompose = WG_SHADER_SOURCE(
//************************************************************************
// compute shaders blend
//************************************************************************
std::string cBlendHeader = WG_SHADER_SOURCE(
@group(0) @binding(0) var imageSrc : texture_storage_2d<rgba8unorm, read>;
@group(0) @binding(1) var imageMsk : texture_storage_2d<rgba8unorm, read>;
@group(0) @binding(2) var imageDst : texture_storage_2d<rgba8unorm, read>;
@group(0) @binding(3) var imageTrg : texture_storage_2d<rgba8unorm, write>;
@group(1) @binding(0) var<uniform> composeMethod : u32;
@group(2) @binding(0) var<uniform> blendMethod : u32;
@group(3) @binding(0) var<uniform> opacity : f32;
@group(1) @binding(0) var imageDst : texture_storage_2d<rgba8unorm, read>;
@group(2) @binding(0) var imageTgt : texture_storage_2d<rgba8unorm, write>;
@group(3) @binding(0) var<uniform> So : 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; };
fn cs_main(@builtin(global_invocation_id) id: vec3u) {
let colorSrc = textureLoad(imageSrc, id.xy);
if (colorSrc.a == 0.0) { return; }
let colorDst = textureLoad(imageDst, id.xy);
var Sc: vec3f = colorSrc.rgb;
var Sa: f32 = colorSrc.a;
let Dc: vec3f = colorDst.rgb;
let Da: f32 = colorDst.a;
var Rc: vec3f = colorDst.rgb;
var Ra: f32 = 1.0;
);
const std::string cBlendPreProcess_Gradient = WG_SHADER_SOURCE(
Sc = Sc + Dc * (1.0 - Sa);
Sa = Sa + Da * (1.0 - Sa);
);
const std::string cBlendPreProcess_Image = WG_SHADER_SOURCE(
Sc = Sc * So;
Sa = Sa * So;
);
const std::string cBlendEquation_Normal = WG_SHADER_SOURCE(
Rc = Sc + Dc * (1.0 - Sa);
Ra = Sa + Da * (1.0 - Sa);
);
const std::string cBlendEquation_Add = WG_SHADER_SOURCE(
Rc = Sc + Dc;
);
const std::string cBlendEquation_Screen = WG_SHADER_SOURCE(
Rc = Sc + Dc - Sc * Dc;
);
const std::string cBlendEquation_Multiply = WG_SHADER_SOURCE(
Rc = Sc * Dc;
);
const std::string cBlendEquation_Overlay = WG_SHADER_SOURCE(
Rc.r = select(1.0 - min(1.0, 2 * (1 - Sc.r) * (1 - Dc.r)), min(1.0, 2 * Sc.r * Dc.r), (Dc.r < 0.5));
Rc.g = select(1.0 - min(1.0, 2 * (1 - Sc.g) * (1 - Dc.g)), min(1.0, 2 * Sc.g * Dc.g), (Dc.g < 0.5));
Rc.b = select(1.0 - min(1.0, 2 * (1 - Sc.b) * (1 - Dc.b)), min(1.0, 2 * Sc.b * Dc.b), (Dc.b < 0.5));
);
const std::string cBlendEquation_Difference = WG_SHADER_SOURCE(
Rc = abs(Dc - Sc);
);
const std::string cBlendEquation_Exclusion = WG_SHADER_SOURCE(
let One = vec3f(1.0, 1.0, 1.0);
Rc = min(One, Sc + Dc - min(One, 2 * Sc * Dc));
);
const std::string cBlendEquation_SrcOver = WG_SHADER_SOURCE(
Rc = Sc; Ra = Sa;
);
const std::string cBlendEquation_Darken = WG_SHADER_SOURCE(
Rc = min(Sc, Dc);
);
const std::string cBlendEquation_Lighten = WG_SHADER_SOURCE(
Rc = max(Sc, Dc);
);
const std::string cBlendEquation_ColorDodge = WG_SHADER_SOURCE(
Rc.r = select(Dc.r, (Dc.r * 255.0 / (255.0 - Sc.r * 255.0))/255.0, (1.0 - Sc.r > 0.0));
Rc.g = select(Dc.g, (Dc.g * 255.0 / (255.0 - Sc.g * 255.0))/255.0, (1.0 - Sc.g > 0.0));
Rc.b = select(Dc.b, (Dc.b * 255.0 / (255.0 - Sc.b * 255.0))/255.0, (1.0 - Sc.b > 0.0));
);
const std::string cBlendEquation_ColorBurn = WG_SHADER_SOURCE(
Rc.r = select(1.0 - Dc.r, (255.0 - (255.0 - Dc.r * 255.0) / (Sc.r * 255.0)) / 255.0, (Sc.r > 0.0));
Rc.g = select(1.0 - Dc.g, (255.0 - (255.0 - Dc.g * 255.0) / (Sc.g * 255.0)) / 255.0, (Sc.g > 0.0));
Rc.b = select(1.0 - Dc.b, (255.0 - (255.0 - Dc.b * 255.0) / (Sc.b * 255.0)) / 255.0, (Sc.b > 0.0));
);
const std::string cBlendEquation_HardLight = WG_SHADER_SOURCE(
Rc.r = select(1.0 - min(1.0, 2 * (1 - Sc.r) * (1 - Dc.r)), min(1.0, 2 * Sc.r * Dc.r), (Sc.r < 0.5));
Rc.g = select(1.0 - min(1.0, 2 * (1 - Sc.g) * (1 - Dc.g)), min(1.0, 2 * Sc.g * Dc.g), (Sc.g < 0.5));
Rc.b = select(1.0 - min(1.0, 2 * (1 - Sc.b) * (1 - Dc.b)), min(1.0, 2 * Sc.b * Dc.b), (Sc.b < 0.5));
);
const std::string cBlendEquation_SoftLight = WG_SHADER_SOURCE(
let One = vec3f(1.0, 1.0, 1.0);
Rc = min(One, (One - 2 * Sc) * Dc * Dc + 2.0 * Sc * Dc);
);
const std::string cBlendPostProcess_Gradient = WG_SHADER_SOURCE(
);
const std::string cBlendPostProcess_Image = WG_SHADER_SOURCE(
Rc = mix(Dc, Rc, Sa);
Ra = mix(Da, Ra, Sa);
);
const std::string cBlendFooter = WG_SHADER_SOURCE(
textureStore(imageTgt, id.xy, vec4f(Rc, Ra));
}
);
std::string blend_solid_Normal = cBlendHeader + cBlendEquation_Normal + cBlendFooter;
std::string blend_solid_Add = cBlendHeader + cBlendEquation_Add + cBlendFooter;
std::string blend_solid_Screen = cBlendHeader + cBlendEquation_Screen + cBlendFooter;
std::string blend_solid_Multiply = cBlendHeader + cBlendEquation_Multiply + cBlendFooter;
std::string blend_solid_Overlay = cBlendHeader + cBlendEquation_Overlay + cBlendFooter;
std::string blend_solid_Difference = cBlendHeader + cBlendEquation_Difference + cBlendFooter;
std::string blend_solid_Exclusion = cBlendHeader + cBlendEquation_Exclusion + cBlendFooter;
std::string blend_solid_SrcOver = cBlendHeader + cBlendEquation_SrcOver + cBlendFooter;
std::string blend_solid_Darken = cBlendHeader + cBlendEquation_Darken + cBlendFooter;
std::string blend_solid_Lighten = cBlendHeader + cBlendEquation_Lighten + cBlendFooter;
std::string blend_solid_ColorDodge = cBlendHeader + cBlendEquation_ColorDodge + cBlendFooter;
std::string blend_solid_ColorBurn = cBlendHeader + cBlendEquation_ColorBurn + cBlendFooter;
std::string blend_solid_HardLight = cBlendHeader + cBlendEquation_HardLight + cBlendFooter;
std::string blend_solid_SoftLight = cBlendHeader + cBlendEquation_SoftLight + cBlendFooter;
const char* cShaderSrc_Blend_Solid[] {
blend_solid_Normal.c_str(),
blend_solid_Add.c_str(),
blend_solid_Screen.c_str(),
blend_solid_Multiply.c_str(),
blend_solid_Overlay.c_str(),
blend_solid_Difference.c_str(),
blend_solid_Exclusion.c_str(),
blend_solid_SrcOver.c_str(),
blend_solid_Darken.c_str(),
blend_solid_Lighten.c_str(),
blend_solid_ColorDodge.c_str(),
blend_solid_ColorBurn.c_str(),
blend_solid_HardLight.c_str(),
blend_solid_SoftLight.c_str()
};
const std::string cBlendHeader_Gradient = cBlendHeader + cBlendPreProcess_Gradient;
const std::string cBlendFooter_Gradient = cBlendPostProcess_Gradient + cBlendFooter;
std::string blend_gradient_Normal = cBlendHeader_Gradient + cBlendEquation_Normal + cBlendFooter_Gradient;
std::string blend_gradient_Add = cBlendHeader_Gradient + cBlendEquation_Add + cBlendFooter_Gradient;
std::string blend_gradient_Screen = cBlendHeader_Gradient + cBlendEquation_Screen + cBlendFooter_Gradient;
std::string blend_gradient_Multiply = cBlendHeader_Gradient + cBlendEquation_Multiply + cBlendFooter_Gradient;
std::string blend_gradient_Overlay = cBlendHeader_Gradient + cBlendEquation_Overlay + cBlendFooter_Gradient;
std::string blend_gradient_Difference = cBlendHeader_Gradient + cBlendEquation_Difference + cBlendFooter_Gradient;
std::string blend_gradient_Exclusion = cBlendHeader_Gradient + cBlendEquation_Exclusion + cBlendFooter_Gradient;
std::string blend_gradient_SrcOver = cBlendHeader_Gradient + cBlendEquation_SrcOver + cBlendFooter_Gradient;
std::string blend_gradient_Darken = cBlendHeader_Gradient + cBlendEquation_Darken + cBlendFooter_Gradient;
std::string blend_gradient_Lighten = cBlendHeader_Gradient + cBlendEquation_Lighten + cBlendFooter_Gradient;
std::string blend_gradient_ColorDodge = cBlendHeader_Gradient + cBlendEquation_ColorDodge + cBlendFooter_Gradient;
std::string blend_gradient_ColorBurn = cBlendHeader_Gradient + cBlendEquation_ColorBurn + cBlendFooter_Gradient;
std::string blend_gradient_HardLight = cBlendHeader_Gradient + cBlendEquation_HardLight + cBlendFooter_Gradient;
std::string blend_gradient_SoftLight = cBlendHeader_Gradient + cBlendEquation_SoftLight + cBlendFooter_Gradient;
const char* cShaderSrc_Blend_Gradient[] {
blend_gradient_Normal.c_str(),
blend_gradient_Add.c_str(),
blend_gradient_Screen.c_str(),
blend_gradient_Multiply.c_str(),
blend_gradient_Overlay.c_str(),
blend_gradient_Difference.c_str(),
blend_gradient_Exclusion.c_str(),
blend_gradient_SrcOver.c_str(),
blend_gradient_Darken.c_str(),
blend_gradient_Lighten.c_str(),
blend_gradient_ColorDodge.c_str(),
blend_gradient_ColorBurn.c_str(),
blend_gradient_HardLight.c_str(),
blend_gradient_SoftLight.c_str()
};
const std::string cBlendHeader_Image = cBlendHeader + cBlendPreProcess_Image;
const std::string cBlendFooter_Image = cBlendPostProcess_Image + cBlendFooter;
std::string blend_image_Normal = cBlendHeader_Image + cBlendEquation_Normal + cBlendFooter;
std::string blend_image_Add = cBlendHeader_Image + cBlendEquation_Add + cBlendFooter_Image;
std::string blend_image_Screen = cBlendHeader_Image + cBlendEquation_Screen + cBlendFooter_Image;
std::string blend_image_Multiply = cBlendHeader_Image + cBlendEquation_Multiply + cBlendFooter_Image;
std::string blend_image_Overlay = cBlendHeader_Image + cBlendEquation_Overlay + cBlendFooter_Image;
std::string blend_image_Difference = cBlendHeader_Image + cBlendEquation_Difference + cBlendFooter_Image;
std::string blend_image_Exclusion = cBlendHeader_Image + cBlendEquation_Exclusion + cBlendFooter_Image;
std::string blend_image_SrcOver = cBlendHeader_Image + cBlendEquation_SrcOver + cBlendFooter_Image;
std::string blend_image_Darken = cBlendHeader_Image + cBlendEquation_Darken + cBlendFooter_Image;
std::string blend_image_Lighten = cBlendHeader_Image + cBlendEquation_Lighten + cBlendFooter_Image;
std::string blend_image_ColorDodge = cBlendHeader_Image + cBlendEquation_ColorDodge + cBlendFooter_Image;
std::string blend_image_ColorBurn = cBlendHeader_Image + cBlendEquation_ColorBurn + cBlendFooter_Image;
std::string blend_image_HardLight = cBlendHeader_Image + cBlendEquation_HardLight + cBlendFooter_Image;
std::string blend_image_SoftLight = cBlendHeader_Image + cBlendEquation_SoftLight + cBlendFooter_Image;
const char* cShaderSrc_Blend_Image[] {
blend_image_Normal.c_str(),
blend_image_Add.c_str(),
blend_image_Screen.c_str(),
blend_image_Multiply.c_str(),
blend_image_Overlay.c_str(),
blend_image_Difference.c_str(),
blend_image_Exclusion.c_str(),
blend_image_SrcOver.c_str(),
blend_image_Darken.c_str(),
blend_image_Lighten.c_str(),
blend_image_ColorDodge.c_str(),
blend_image_ColorBurn.c_str(),
blend_image_HardLight.c_str(),
blend_image_SoftLight.c_str()
};
//************************************************************************
// compute shaders compose
//************************************************************************
std::string cComposeHeader = WG_SHADER_SOURCE(
@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>;
@compute @workgroup_size(8, 8)
fn cs_main(@builtin(global_invocation_id) id: vec3u) {
let colorSrc = textureLoad(imageSrc, id.xy);
let colorMsk = textureLoad(imageMsk, id.xy);
let colorDst = textureLoad(imageDst, id.xy);
var One: vec3f = vec3(1.0);
let luma: f32 = dot(colorMsk.xyz, vec3f(0.2125, 0.7154, 0.0721));
var So : f32 = opacity;
var Mc : vec3f = colorMsk.rgb;
var Ma : f32 = colorMsk.a;
var MSc : vec3f = colorSrc.rgb;
var MSa : f32 = colorSrc.a;
var Sc : vec3f = colorSrc.rgb;
var Sa : f32 = colorSrc.a;
var Dc : vec3f = colorDst.rgb;
var Da : f32 = colorDst.a;
var Rc : vec3f = colorDst.rgb;
var Ra : f32 = 1.0;
var Sc : vec3f = colorSrc.rgb;
var Sa : f32 = colorSrc.a;
var Mc : vec3f = colorMsk.rgb;
var Ma : f32 = colorMsk.a;
var Dc : vec3f = colorDst.rgb;
var Da : f32 = colorDst.a;
var Rc : vec3f = colorDst.rgb;
var Ra : f32 = colorDst.a;
);
if (composeMethod <= /* InvLumaMask */5u) {
switch composeMethod {
/* AlphaMask */ case 2u: { Sc = MSc * Ma; Sa = MSa * Ma; }
/* InvAlphaMask */ case 3u: { Sc = MSc * (1.0 - Ma); Sa = MSa * (1.0 - Ma); }
/* LumaMask */ case 4u: { Sc = MSc * luma; Sa = MSa * luma; }
/* InvLumaMask */ case 5u: { Sc = MSc * (1.0-luma); Sa = MSa * (1.0-luma); }
default: { Sc = MSc; Sa = MSa; }
}
Rc = Sc + Dc * (1.0 - Sa);
Ra = Sa + Da * (1.0 - Sa);
} else {
Sc = Dc;
switch composeMethod {
/* AddMask */ case 6u: { Sa = MSa + Ma * (1.0 - MSa); }
/* SubtractMask */ case 7u: { Sa = MSa * (1.0 - Ma); }
/* IntersectMask */ case 8u: { Sa = MSa * Ma; }
/* DifferenceMask */ case 9u: { Sa = MSa * (1.0 - Ma) + Ma * (1.0 - MSa); }
default: { Rc = Sc; Ra = Sa; }
}
Rc = Sc;
Ra = Sa;
}
std::string cComposeEquation_None = WG_SHADER_SOURCE(
Rc = Dc;
Ra = Da;
);
std::string cComposeEquation_ClipPath = WG_SHADER_SOURCE(
Rc = Sc * Ma + Dc * (1.0 - Sa * Ma);
Ra = Sa * Ma + Da * (1.0 - Sa * Ma);
);
textureStore(imageTrg, id.xy, vec4f(Rc, Ra));
std::string cComposeEquation_AlphaMask = WG_SHADER_SOURCE(
Rc = Sc * Ma + Dc * (1.0 - Sa * Ma);
Ra = Sa * Ma + Da * (1.0 - Sa * Ma);
);
std::string cComposeEquation_InvAlphaMask = WG_SHADER_SOURCE(
Rc = Sc * (1.0 - Ma) + Dc * (1.0 - Sa * (1.0 - Ma));
Ra = Sa * (1.0 - Ma) + Da * (1.0 - Sa * (1.0 - Ma));
);
std::string cComposeEquation_LumaMask = WG_SHADER_SOURCE(
let luma: f32 = dot(Mc, vec3f(0.2125, 0.7154, 0.0721));
Rc = Sc * luma + Dc * (1.0 - Sa * luma);
Ra = Sa * luma + Da * (1.0 - Sa * luma);
);
std::string cComposeEquation_InvLumaMask = WG_SHADER_SOURCE(
let luma: f32 = dot(Mc, vec3f(0.2125, 0.7154, 0.0721));
Rc = Sc * (1.0 - luma) + Dc * (1.0 - Sa * (1.0 - luma));
Ra = Sa * (1.0 - luma) + Da * (1.0 - Sa * (1.0 - luma));
);
std::string cComposeEquation_AddMask = WG_SHADER_SOURCE(
Rc = Sc;
Ra = Sa + Ma * (1.0 - Sa);
);
std::string cComposeEquation_SubtractMask = WG_SHADER_SOURCE(
Rc = Sc;
Ra = Sa * (1.0 - Ma);
);
std::string cComposeEquation_IntersectMask = WG_SHADER_SOURCE(
Rc = Sc;
Ra = Sa * Ma;
);
std::string cComposeEquation_DifferenceMask = WG_SHADER_SOURCE(
Rc = Sc;
Ra = Sa * (1.0 - Ma) + Ma * (1.0 - Sa);
);
std::string cComposeFooter = WG_SHADER_SOURCE(
textureStore(imageTgt, id.xy, vec4f(Rc, Ra));
}
);
// pipeline shader modules anti-aliasing
const char* cShaderSource_PipelineComputeAntiAlias = WG_SHADER_SOURCE(
std::string compose_None = cComposeHeader + cComposeEquation_None + cComposeFooter;
std::string compose_ClipPath = cComposeHeader + cComposeEquation_ClipPath + cComposeFooter;
std::string compose_AlphaMask = cComposeHeader + cComposeEquation_AlphaMask + cComposeFooter;
std::string compose_InvAlphaMask = cComposeHeader + cComposeEquation_InvAlphaMask + cComposeFooter;
std::string compose_LumaMask = cComposeHeader + cComposeEquation_LumaMask + cComposeFooter;
std::string compose_InvLumaMask = cComposeHeader + cComposeEquation_InvLumaMask + cComposeFooter;
std::string compose_AddMask = cComposeHeader + cComposeEquation_AddMask + cComposeFooter;
std::string compose_SubtractMask = cComposeHeader + cComposeEquation_SubtractMask + cComposeFooter;
std::string compose_IntersectMask = cComposeHeader + cComposeEquation_IntersectMask + cComposeFooter;
std::string compose_DifferenceMask = cComposeHeader + cComposeEquation_DifferenceMask + cComposeFooter;
const char* cShaderSrc_Compose[10] {
compose_None.c_str(),
compose_ClipPath.c_str(),
compose_AlphaMask.c_str(),
compose_InvAlphaMask.c_str(),
compose_LumaMask.c_str(),
compose_InvLumaMask.c_str(),
compose_AddMask.c_str(),
compose_SubtractMask.c_str(),
compose_IntersectMask.c_str(),
compose_DifferenceMask.c_str(),
};
const char* cShaderSrc_Copy = WG_SHADER_SOURCE(
@group(0) @binding(0) var imageSrc : texture_storage_2d<rgba8unorm, read>;
@group(1) @binding(0) var imageDst : texture_storage_2d<bgra8unorm, write>;
@compute @workgroup_size(8, 8)
fn cs_main( @builtin(global_invocation_id) id: vec3u) {
let texSizeSrc = textureDimensions(imageSrc);
let texSizeDst = textureDimensions(imageDst);
if ((id.x >= texSizeDst.x) || (id.y >= texSizeDst.y)) { return; };
let samples = u32(texSizeSrc.x / texSizeDst.x);
var color = vec4f(0);
for (var i = 0u; i < samples; i++) {
for (var j = 0u; j < samples; j++) {
color += textureLoad(imageSrc, vec2(id.x * samples + j, id.y * samples + i));
}
}
textureStore(imageDst, id.xy, color / f32(samples * samples));
fn cs_main(@builtin(global_invocation_id) id: vec3u) {
textureStore(imageDst, id.xy, textureLoad(imageSrc, id.xy));
}
);

44
src/renderer/wg_engine/tvgWgShaderSrc.h Normal file → Executable file
View file

@ -20,37 +20,23 @@
* SOFTWARE.
*/
#pragma once
#ifndef _TVG_WG_SHADER_SRC_H_
#define _TVG_WG_SHADER_SRC_H_
//*****************************************************************************
// render shader modules
//*****************************************************************************
// graphics shader sources
extern const char* cShaderSrc_Stencil;
extern const char* cShaderSrc_Solid;
extern const char* cShaderSrc_Linear;
extern const char* cShaderSrc_Radial;
extern const char* cShaderSrc_Image;
// pipeline shader modules fill
extern const char* cShaderSource_PipelineFill;
extern const char* cShaderSource_PipelineSolid;
extern const char* cShaderSource_PipelineLinear;
extern const char* cShaderSource_PipelineRadial;
extern const char* cShaderSource_PipelineImage;
// compute shader sources: blend, compose and merge path
extern const char* cShaderSrc_MergeMasks;
extern const char* cShaderSrc_Blend_Solid[14];
extern const char* cShaderSrc_Blend_Solid[14];
extern const char* cShaderSrc_Blend_Gradient[14];
extern const char* cShaderSrc_Blend_Image[14];
extern const char* cShaderSrc_Compose[10];
extern const char* cShaderSrc_Copy;
//*****************************************************************************
// compute shader modules
//*****************************************************************************
// pipeline shader modules clear, compose and blend
extern const char* cShaderSource_PipelineComputeClear;
extern const char* cShaderSource_PipelineComputeCopy;
extern const char* cShaderSource_PipelineComputeBlendSolid;
extern const char* cShaderSource_PipelineComputeBlendGradient;
extern const char* cShaderSource_PipelineComputeBlendImage;
extern const char* cShaderSource_PipelineComputeBlendSolidMask;
extern const char* cShaderSource_PipelineComputeBlendGradientMask;
extern const char* cShaderSource_PipelineComputeBlendImageMask;
extern const char* cShaderSource_PipelineComputeMaskCompose;
extern const char* cShaderSource_PipelineComputeCompose;
extern const char* cShaderSource_PipelineComputeAntiAlias;
#endif // _TVG_WG_SHADER_SRC_H_
#endif // _TVG_WG_SHEDER_SRC_H_

0
src/renderer/wg_engine/tvgWgShaderTypes.cpp Normal file → Executable file
View file

2
src/renderer/wg_engine/tvgWgShaderTypes.h Normal file → Executable file
View file

@ -23,7 +23,7 @@
#ifndef _TVG_WG_SHADER_TYPES_H_
#define _TVG_WG_SHADER_TYPES_H_
#include "tvgWgCommon.h"
#include "tvgRender.h"
///////////////////////////////////////////////////////////////////////////////
// shader types