From 9efab30b0095bff4f5d64b0f0433969cdf44d5fc Mon Sep 17 00:00:00 2001 From: Sergii Liebodkin Date: Sun, 11 Aug 2024 21:24:46 +0000 Subject: [PATCH] webgpu: shaders refactoring Deep shader refactoring for the following purposes: * used pre-calculated gradient texture instead of per-pixel gradient map computation * used HW wrap samples for fill spread setting * unified gradient shader types * used single shader module for composition instead of signle module per composition type * used single shader module for blending for each of fill type (solid, gradient, image) instaed of signle module per blend type * much easier add new composition and blend equations * get rided std::string uasge * shaders code is more readable --- src/renderer/wg_engine/tvgWgBindGroups.cpp | 24 + src/renderer/wg_engine/tvgWgBindGroups.h | 2 + src/renderer/wg_engine/tvgWgCommon.cpp | 46 +- src/renderer/wg_engine/tvgWgCommon.h | 9 +- src/renderer/wg_engine/tvgWgCompositor.cpp | 8 +- src/renderer/wg_engine/tvgWgPipelines.cpp | 122 ++- src/renderer/wg_engine/tvgWgPipelines.h | 14 +- src/renderer/wg_engine/tvgWgRenderData.cpp | 39 +- src/renderer/wg_engine/tvgWgRenderData.h | 8 +- src/renderer/wg_engine/tvgWgRenderTarget.cpp | 2 +- src/renderer/wg_engine/tvgWgRenderer.cpp | 3 +- src/renderer/wg_engine/tvgWgShaderSrc.cpp | 902 ++++++++----------- src/renderer/wg_engine/tvgWgShaderSrc.h | 10 +- src/renderer/wg_engine/tvgWgShaderTypes.cpp | 89 +- src/renderer/wg_engine/tvgWgShaderTypes.h | 70 +- 15 files changed, 633 insertions(+), 715 deletions(-) diff --git a/src/renderer/wg_engine/tvgWgBindGroups.cpp b/src/renderer/wg_engine/tvgWgBindGroups.cpp index 012fe4ee..4478d8d2 100755 --- a/src/renderer/wg_engine/tvgWgBindGroups.cpp +++ b/src/renderer/wg_engine/tvgWgBindGroups.cpp @@ -33,6 +33,18 @@ WGPUBindGroup WgBindGroupLayouts::createBindGroupTexSampled(WGPUSampler sampler, } +WGPUBindGroup WgBindGroupLayouts::createBindGroupTexSampledBuff1Un(WGPUSampler sampler, WGPUTextureView texView, WGPUBuffer buff) +{ + const WGPUBindGroupEntry bindGroupEntrys[] = { + { .binding = 0, .sampler = sampler }, + { .binding = 1, .textureView = texView }, + { .binding = 2, .buffer = buff, .size = wgpuBufferGetSize(buff) } + }; + const WGPUBindGroupDescriptor bindGroupDesc { .layout = layoutTexSampledBuff1Un, .entryCount = 3, .entries = bindGroupEntrys }; + return wgpuDeviceCreateBindGroup(device, &bindGroupDesc); +} + + WGPUBindGroup WgBindGroupLayouts::createBindGroupScreen1WO(WGPUTextureView texView) { const WGPUBindGroupEntry bindGroupEntrys[] = { { .binding = 0, .textureView = texView } @@ -151,6 +163,17 @@ void WgBindGroupLayouts::initialize(WgContext& context) assert(layoutTexSampled); } + { // bind group layout tex sampled with buffer uniform + const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { + { .binding = 0, .visibility = visibility_frag, .sampler = sampler }, + { .binding = 1, .visibility = visibility_frag, .texture = texture }, + { .binding = 2, .visibility = visibility_vert, .buffer = bufferUniform } + }; + const WGPUBindGroupLayoutDescriptor bindGroupLayoutDesc { .entryCount = 3, .entries = bindGroupLayoutEntries }; + layoutTexSampledBuff1Un = wgpuDeviceCreateBindGroupLayout(device, &bindGroupLayoutDesc); + assert(layoutTexSampledBuff1Un); + } + { // bind group layout tex screen 1 RO const WGPUBindGroupLayoutEntry bindGroupLayoutEntries[] { { .binding = 0, .visibility = visibility_frag, .storageTexture = storageScreenWO } @@ -241,6 +264,7 @@ void WgBindGroupLayouts::release(WgContext& context) wgpuBindGroupLayoutRelease(layoutTexStrorage2RO); wgpuBindGroupLayoutRelease(layoutTexStrorage1RO); wgpuBindGroupLayoutRelease(layoutTexStrorage1WO); + wgpuBindGroupLayoutRelease(layoutTexSampledBuff1Un); wgpuBindGroupLayoutRelease(layoutTexSampled); device = nullptr; } diff --git a/src/renderer/wg_engine/tvgWgBindGroups.h b/src/renderer/wg_engine/tvgWgBindGroups.h index 6780442e..88588dec 100755 --- a/src/renderer/wg_engine/tvgWgBindGroups.h +++ b/src/renderer/wg_engine/tvgWgBindGroups.h @@ -30,6 +30,7 @@ private: WGPUDevice device{}; public: WGPUBindGroupLayout layoutTexSampled{}; + WGPUBindGroupLayout layoutTexSampledBuff1Un{}; WGPUBindGroupLayout layoutTexScreen1WO{}; WGPUBindGroupLayout layoutTexStrorage1WO{}; WGPUBindGroupLayout layoutTexStrorage1RO{}; @@ -40,6 +41,7 @@ public: WGPUBindGroupLayout layoutBuffer3Un{}; public: WGPUBindGroup createBindGroupTexSampled(WGPUSampler sampler, WGPUTextureView texView); + WGPUBindGroup createBindGroupTexSampledBuff1Un(WGPUSampler sampler, WGPUTextureView texView, WGPUBuffer buff); WGPUBindGroup createBindGroupScreen1WO(WGPUTextureView texView); WGPUBindGroup createBindGroupStrorage1WO(WGPUTextureView texView); WGPUBindGroup createBindGroupStrorage1RO(WGPUTextureView texView); diff --git a/src/renderer/wg_engine/tvgWgCommon.cpp b/src/renderer/wg_engine/tvgWgCommon.cpp index ad43f54c..c97cf703 100755 --- a/src/renderer/wg_engine/tvgWgCommon.cpp +++ b/src/renderer/wg_engine/tvgWgCommon.cpp @@ -69,26 +69,33 @@ void WgContext::initialize(WGPUInstance instance, WGPUSurface surface) // create shared webgpu assets allocateBufferIndexFan(32768); - samplerNearest = createSampler(WGPUFilterMode_Nearest, WGPUMipmapFilterMode_Nearest); - samplerLinear = createSampler(WGPUFilterMode_Linear, WGPUMipmapFilterMode_Linear); - assert(samplerNearest); - assert(samplerLinear); + samplerNearestRepeat = createSampler(WGPUFilterMode_Nearest, WGPUMipmapFilterMode_Nearest, WGPUAddressMode_Repeat); + samplerLinearRepeat = createSampler(WGPUFilterMode_Linear, WGPUMipmapFilterMode_Linear, WGPUAddressMode_Repeat); + samplerLinearMirror = createSampler(WGPUFilterMode_Linear, WGPUMipmapFilterMode_Linear, WGPUAddressMode_MirrorRepeat); + samplerLinearClamp = createSampler(WGPUFilterMode_Linear, WGPUMipmapFilterMode_Linear, WGPUAddressMode_ClampToEdge); + assert(samplerNearestRepeat); + assert(samplerLinearRepeat); + assert(samplerLinearMirror); + assert(samplerLinearClamp); } void WgContext::release() { - releaseSampler(samplerNearest); - releaseSampler(samplerLinear); + releaseSampler(samplerLinearClamp); + releaseSampler(samplerLinearMirror); + releaseSampler(samplerLinearRepeat); + releaseSampler(samplerNearestRepeat); releaseBuffer(bufferIndexFan); wgpuDeviceRelease(device); wgpuAdapterRelease(adapter); } -WGPUSampler WgContext::createSampler(WGPUFilterMode filter, WGPUMipmapFilterMode mipmapFilter) +WGPUSampler WgContext::createSampler(WGPUFilterMode filter, WGPUMipmapFilterMode mipmapFilter, WGPUAddressMode addrMode) { const WGPUSamplerDescriptor samplerDesc { + .addressModeU = addrMode, .addressModeV = addrMode, .addressModeW = addrMode, .magFilter = filter, .minFilter = filter, .mipmapFilter = mipmapFilter, .lodMinClamp = 0.0f, .lodMaxClamp = 32.0f, .maxAnisotropy = 1 }; @@ -96,6 +103,31 @@ WGPUSampler WgContext::createSampler(WGPUFilterMode filter, WGPUMipmapFilterMode } +bool WgContext::allocateTexture(WGPUTexture& texture, uint32_t width, uint32_t height, WGPUTextureFormat format, void* data) +{ + if ((texture) && (wgpuTextureGetWidth(texture) == width) && (wgpuTextureGetHeight(texture) == height)) { + // update texture data + const WGPUImageCopyTexture imageCopyTexture{ .texture = texture }; + const WGPUTextureDataLayout textureDataLayout{ .bytesPerRow = 4 * width, .rowsPerImage = height }; + const WGPUExtent3D writeSize{ .width = width, .height = height, .depthOrArrayLayers = 1 }; + wgpuQueueWriteTexture(queue, &imageCopyTexture, data, 4 * width * height, &textureDataLayout, &writeSize); + wgpuQueueSubmit(queue, 0, nullptr); + } else { + releaseTexture(texture); + texture = createTexture(width, height, format); + // update texture data + const WGPUImageCopyTexture imageCopyTexture{ .texture = texture }; + const WGPUTextureDataLayout textureDataLayout{ .bytesPerRow = 4 * width, .rowsPerImage = height }; + const WGPUExtent3D writeSize{ .width = width, .height = height, .depthOrArrayLayers = 1 }; + wgpuQueueWriteTexture(queue, &imageCopyTexture, data, 4 * width * height, &textureDataLayout, &writeSize); + wgpuQueueSubmit(queue, 0, nullptr); + return true; + } + return false; + +} + + WGPUTexture WgContext::createTexture(uint32_t width, uint32_t height, WGPUTextureFormat format) { const WGPUTextureDescriptor textureDesc { diff --git a/src/renderer/wg_engine/tvgWgCommon.h b/src/renderer/wg_engine/tvgWgCommon.h index b6095caa..a73fd256 100755 --- a/src/renderer/wg_engine/tvgWgCommon.h +++ b/src/renderer/wg_engine/tvgWgCommon.h @@ -44,18 +44,21 @@ struct WgContext { WgPipelines* pipelines{}; // shared webgpu assets WGPUBuffer bufferIndexFan{}; - WGPUSampler samplerNearest{}; - WGPUSampler samplerLinear{}; + WGPUSampler samplerNearestRepeat{}; + WGPUSampler samplerLinearRepeat{}; + WGPUSampler samplerLinearMirror{}; + WGPUSampler samplerLinearClamp{}; void initialize(WGPUInstance instance, WGPUSurface surface); void release(); // create common objects - WGPUSampler createSampler(WGPUFilterMode filter, WGPUMipmapFilterMode mipmapFilter); + WGPUSampler createSampler(WGPUFilterMode filter, WGPUMipmapFilterMode mipmapFilter, WGPUAddressMode addrMode); 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); + bool allocateTexture(WGPUTexture& texture, uint32_t width, uint32_t height, WGPUTextureFormat format, void* data); // release common objects void releaseTextureView(WGPUTextureView& textureView); diff --git a/src/renderer/wg_engine/tvgWgCompositor.cpp b/src/renderer/wg_engine/tvgWgCompositor.cpp index a1d5368d..7b639ee3 100755 --- a/src/renderer/wg_engine/tvgWgCompositor.cpp +++ b/src/renderer/wg_engine/tvgWgCompositor.cpp @@ -371,10 +371,10 @@ void WgCompositor::drawShape(WgContext& context, WgRenderDataShape* renderData, 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); + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 2, settings.bindGroupGradient, 0, nullptr); wgpuRenderPassEncoderSetPipeline(renderPassEncoder, pipelines->linear[blendTypeInd]); } else if (settings.fillType == WgRenderSettingsType::Radial) { - wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 2, settings.bindGroupRadial, 0, nullptr); + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 2, settings.bindGroupGradient, 0, nullptr); wgpuRenderPassEncoderSetPipeline(renderPassEncoder, pipelines->radial[blendTypeInd]); } // draw to color (second pass) @@ -409,10 +409,10 @@ void WgCompositor::drawStrokes(WgContext& context, WgRenderDataShape* renderData 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); + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 2, settings.bindGroupGradient, 0, nullptr); wgpuRenderPassEncoderSetPipeline(renderPassEncoder, pipelines->linear[blendTypeInd]); } else if (settings.fillType == WgRenderSettingsType::Radial) { - wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 2, settings.bindGroupRadial, 0, nullptr); + wgpuRenderPassEncoderSetBindGroup(renderPassEncoder, 2, settings.bindGroupGradient, 0, nullptr); wgpuRenderPassEncoderSetPipeline(renderPassEncoder, pipelines->radial[blendTypeInd]); } // draw to color (second pass) diff --git a/src/renderer/wg_engine/tvgWgPipelines.cpp b/src/renderer/wg_engine/tvgWgPipelines.cpp index 7709c0df..465844bb 100755 --- a/src/renderer/wg_engine/tvgWgPipelines.cpp +++ b/src/renderer/wg_engine/tvgWgPipelines.cpp @@ -22,6 +22,7 @@ #include "tvgWgPipelines.h" #include "tvgWgShaderSrc.h" +#include WGPUShaderModule WgPipelines::createShaderModule(WGPUDevice device, const char* label, const char* code) { @@ -74,14 +75,15 @@ WGPURenderPipeline WgPipelines::createRenderPipeline( WGPUComputePipeline WgPipelines::createComputePipeline( WGPUDevice device, const char* pipelineLabel, - const WGPUShaderModule shaderModule, const WGPUPipelineLayout pipelineLayout) + const WGPUShaderModule shaderModule, const char* entryPoint, + const WGPUPipelineLayout pipelineLayout) { const WGPUComputePipelineDescriptor computePipelineDesc{ .label = pipelineLabel, .layout = pipelineLayout, .compute = { .module = shaderModule, - .entryPoint = "cs_main" + .entryPoint = entryPoint } }; return wgpuDeviceCreateComputePipeline(device, &computePipelineDesc); @@ -165,11 +167,16 @@ void WgPipelines::initialize(WgContext& context) layouts.layoutBuffer1Un, layouts.layoutBuffer2Un }; - const WGPUBindGroupLayout bindGroupLayoutsFill[] = { + const WGPUBindGroupLayout bindGroupLayoutsSolid[] = { layouts.layoutBuffer1Un, layouts.layoutBuffer2Un, layouts.layoutBuffer1Un }; + const WGPUBindGroupLayout bindGroupLayoutsGradient[] = { + layouts.layoutBuffer1Un, + layouts.layoutBuffer2Un, + layouts.layoutTexSampledBuff1Un + }; const WGPUBindGroupLayout bindGroupLayoutsImage[] = { layouts.layoutBuffer1Un, layouts.layoutBuffer2Un, @@ -199,7 +206,8 @@ void WgPipelines::initialize(WgContext& context) // pipeline layouts layoutStencil = createPipelineLayout(context.device, bindGroupLayoutsStencil, 2); - layoutFill = createPipelineLayout(context.device, bindGroupLayoutsFill, 3); + layoutSolid = createPipelineLayout(context.device, bindGroupLayoutsSolid, 3); + layoutGradient = createPipelineLayout(context.device, bindGroupLayoutsGradient, 3); layoutImage = createPipelineLayout(context.device, bindGroupLayoutsImage, 3); layoutBlend = createPipelineLayout(context.device, bindGroupLayoutsBlend, 4); layoutCompose = createPipelineLayout(context.device, bindGroupLayoutsCompose, 4); @@ -257,7 +265,7 @@ void WgPipelines::initialize(WgContext& context) for (uint32_t i = 0; i < 3; i++) { solid[i] = createRenderPipeline( context.device, "The render pipeline solid", - shaderSolid, layoutFill, + shaderSolid, layoutSolid, vertexBufferLayoutsShape, 1, WGPUColorWriteMask_All, WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero, @@ -269,7 +277,7 @@ void WgPipelines::initialize(WgContext& context) for (uint32_t i = 0; i < 3; i++) { radial[i] = createRenderPipeline( context.device, "The render pipeline radial", - shaderRadial, layoutFill, + shaderRadial, layoutGradient, vertexBufferLayoutsShape, 1, WGPUColorWriteMask_All, WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero, @@ -281,7 +289,7 @@ void WgPipelines::initialize(WgContext& context) for (uint32_t i = 0; i < 3; i++) { linear[i] = createRenderPipeline( context.device, "The render pipeline linear", - shaderLinear, layoutFill, + shaderLinear, layoutGradient, vertexBufferLayoutsShape, 1, WGPUColorWriteMask_All, WGPUCompareFunction_NotEqual, WGPUStencilOperation_Zero, @@ -302,26 +310,68 @@ void WgPipelines::initialize(WgContext& context) } // 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); + mergeMasks = createComputePipeline(context.device, "The pipeline merge masks", shaderMergeMasks, "cs_main", layoutMergeMasks); + copy = createComputePipeline(context.device, "The pipeline copy", shaderCopy, "cs_main", layoutCopy); + + // compute shader blend names + const char* shaderBlendNames[] { + "cs_main_Normal", + "cs_main_Add", + "cs_main_Screen", + "cs_main_Multiply", + "cs_main_Overlay", + "cs_main_Difference", + "cs_main_Exclusion", + "cs_main_SrcOver", + "cs_main_Darken", + "cs_main_Lighten", + "cs_main_ColorDodge", + "cs_main_ColorBurn", + "cs_main_HardLight", + "cs_main_SoftLight" + }; + + // create blend shaders + char shaderSourceBuff[16384]; + shaderBlendSolid = createShaderModule(context.device, "The shader blend solid", strcat(strcpy(shaderSourceBuff, cShaderSrc_BlendHeader_Solid), cShaderSrc_Blend_Funcs)); + shaderBlendGradient = createShaderModule(context.device, "The shader blend gradient", strcat(strcpy(shaderSourceBuff, cShaderSrc_BlendHeader_Gradient), cShaderSrc_Blend_Funcs)); + shaderBlendImage = createShaderModule(context.device, "The shader blend image", strcat(strcpy(shaderSourceBuff, cShaderSrc_BlendHeader_Image), cShaderSrc_Blend_Funcs)); + + // create blend pipelines + const size_t shaderBlendNamesCnt = sizeof(shaderBlendNames) / sizeof(shaderBlendNames[0]); + const size_t pipesBlendCnt = sizeof(blendSolid) / sizeof(blendSolid[0]); + assert(shaderBlendNamesCnt == pipesBlendCnt); + for (uint32_t i = 0; i < pipesBlendCnt; i++) { + blendSolid[i] = createComputePipeline(context.device, "The pipeline blend solid", shaderBlendSolid, shaderBlendNames[i], layoutBlend); + blendGradient[i] = createComputePipeline(context.device, "The pipeline blend gradient", shaderBlendGradient, shaderBlendNames[i], layoutBlend); + blendImage[i] = createComputePipeline(context.device, "The pipeline blend image", shaderBlendImage, shaderBlendNames[i], layoutBlend); } + + // compute shader compose names + const char* shaderComposeNames[] { + "cs_main_None", + "cs_main_ClipPath", + "cs_main_AlphaMask", + "cs_main_InvAlphaMask", + "cs_main_LumaMask", + "cs_main_InvLumaMask", + "cs_main_AddMask", + "cs_main_SubtractMask", + "cs_main_IntersectMask", + "cs_main_DifferenceMask", + "cs_main_LightenMask", + "cs_main_DarkenMask" + }; + + // create compose shaders + shaderCompose = createShaderModule(context.device, "The shader compose", cShaderSrc_Compose); + + // create compose pipelines + const size_t shaderComposeNamesCnt = sizeof(shaderComposeNames) / sizeof(shaderComposeNames[0]); + const size_t pipesComposeCnt = sizeof(compose) / sizeof(compose[0]); + assert(shaderComposeNamesCnt == pipesComposeCnt); + for (uint32_t i = 0; i < pipesComposeCnt; i++) + compose[i] = createComputePipeline(context.device, "The pipeline compose", shaderCompose, shaderComposeNames[i], layoutCompose); } void WgPipelines::releaseGraphicHandles(WgContext& context) @@ -337,7 +387,8 @@ void WgPipelines::releaseGraphicHandles(WgContext& context) releaseRenderPipeline(evenodd); releaseRenderPipeline(winding); releasePipelineLayout(layoutImage); - releasePipelineLayout(layoutFill); + releasePipelineLayout(layoutGradient); + releasePipelineLayout(layoutSolid); releasePipelineLayout(layoutStencil); releaseShaderModule(shaderImage); releaseShaderModule(shaderLinear); @@ -349,19 +400,14 @@ void WgPipelines::releaseGraphicHandles(WgContext& context) void WgPipelines::releaseComputeHandles(WgContext& context) { - const size_t shaderComposeCnt = sizeof(shaderCompose)/sizeof(shaderCompose[0]); - for (uint32_t i = 0; i < shaderComposeCnt; i++) { + size_t pipesComposeCnt = sizeof(compose) / sizeof(compose[0]); + for (uint32_t i = 0; i < pipesComposeCnt; i++) releaseComputePipeline(compose[i]); - releaseShaderModule(shaderCompose[i]); - } - const size_t shaderBlendImageCnt = sizeof(shaderBlendImage)/sizeof(shaderBlendImage[0]); - for (uint32_t i = 0; i < shaderBlendImageCnt; i++) { + const size_t pipesBlendCnt = sizeof(blendSolid)/sizeof(blendSolid[0]); + for (uint32_t i = 0; i < pipesBlendCnt; i++) { releaseComputePipeline(blendImage[i]); releaseComputePipeline(blendSolid[i]); releaseComputePipeline(blendGradient[i]); - releaseShaderModule(shaderBlendImage[i]); - releaseShaderModule(shaderBlendGradient[i]); - releaseShaderModule(shaderBlendSolid[i]); } releaseComputePipeline(copy); releaseComputePipeline(mergeMasks); @@ -369,6 +415,10 @@ void WgPipelines::releaseComputeHandles(WgContext& context) releasePipelineLayout(layoutBlend); releasePipelineLayout(layoutMergeMasks); releasePipelineLayout(layoutCopy); + releaseShaderModule(shaderCompose); + releaseShaderModule(shaderBlendImage); + releaseShaderModule(shaderBlendGradient); + releaseShaderModule(shaderBlendSolid); releaseShaderModule(shaderMergeMasks); releaseShaderModule(shaderCopy); } diff --git a/src/renderer/wg_engine/tvgWgPipelines.h b/src/renderer/wg_engine/tvgWgPipelines.h index 4e209ac5..4cb84d6c 100755 --- a/src/renderer/wg_engine/tvgWgPipelines.h +++ b/src/renderer/wg_engine/tvgWgPipelines.h @@ -37,15 +37,16 @@ private: WGPUShaderModule shaderImage{}; // compute pipeline shaders WGPUShaderModule shaderMergeMasks; - WGPUShaderModule shaderBlendSolid[14]; - WGPUShaderModule shaderBlendGradient[14]; - WGPUShaderModule shaderBlendImage[14]; - WGPUShaderModule shaderCompose[12]; + WGPUShaderModule shaderBlendSolid; + WGPUShaderModule shaderBlendGradient; + WGPUShaderModule shaderBlendImage; + WGPUShaderModule shaderCompose; WGPUShaderModule shaderCopy; private: // graphics pipeline layouts WGPUPipelineLayout layoutStencil{}; - WGPUPipelineLayout layoutFill{}; + WGPUPipelineLayout layoutSolid{}; + WGPUPipelineLayout layoutGradient{}; WGPUPipelineLayout layoutImage{}; // compute pipeline layouts WGPUPipelineLayout layoutMergeMasks{}; @@ -86,7 +87,8 @@ private: const WGPUPrimitiveState primitiveState, const WGPUMultisampleState multisampleState, const WGPUBlendState blendState); WGPUComputePipeline createComputePipeline( WGPUDevice device, const char* pipelineLabel, - const WGPUShaderModule shaderModule, const WGPUPipelineLayout pipelineLayout); + const WGPUShaderModule shaderModule, const char* entryPoint, + const WGPUPipelineLayout pipelineLayout); void releaseComputePipeline(WGPUComputePipeline& computePipeline); void releaseRenderPipeline(WGPURenderPipeline& renderPipeline); void releasePipelineLayout(WGPUPipelineLayout& pipelineLayout); diff --git a/src/renderer/wg_engine/tvgWgRenderData.cpp b/src/renderer/wg_engine/tvgWgRenderData.cpp index 9e7c9676..3ebee08a 100755 --- a/src/renderer/wg_engine/tvgWgRenderData.cpp +++ b/src/renderer/wg_engine/tvgWgRenderData.cpp @@ -212,22 +212,30 @@ 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 + WgShaderTypeGradient gradient; if (fill->type() == Type::LinearGradient) { - WgShaderTypeLinearGradient linearGradient((LinearGradient*)fill); - if (context.allocateBufferUniform(bufferGroupLinear, &linearGradient, sizeof(linearGradient))) { - context.pipelines->layouts.releaseBindGroup(bindGroupLinear); - bindGroupLinear = context.pipelines->layouts.createBindGroupBuffer1Un(bufferGroupLinear); - } + gradient.update((LinearGradient*)fill); fillType = WgRenderSettingsType::Linear; } else if (fill->type() == Type::RadialGradient) { - WgShaderTypeRadialGradient radialGradient((RadialGradient*)fill); - if (context.allocateBufferUniform(bufferGroupRadial, &radialGradient, sizeof(radialGradient))) { - context.pipelines->layouts.releaseBindGroup(bindGroupRadial); - bindGroupRadial = context.pipelines->layouts.createBindGroupBuffer1Un(bufferGroupRadial); - } + gradient.update((RadialGradient*)fill); fillType = WgRenderSettingsType::Radial; } + // update gpu assets + bool bufferGradientSettingsChanged = context.allocateBufferUniform(bufferGroupGradient, &gradient.settings, sizeof(gradient.settings)); + bool textureGradientChanged = context.allocateTexture(texGradient, WG_TEXTURE_GRADIENT_SIZE, 1, WGPUTextureFormat_RGBA8Unorm, gradient.texData); + if (bufferGradientSettingsChanged || textureGradientChanged) { + // update texture view + context.releaseTextureView(texViewGradient); + texViewGradient = context.createTextureView(texGradient); + // get sampler by spread type + WGPUSampler sampler = context.samplerLinearClamp; + if (fill->spread() == FillSpread::Reflect) sampler = context.samplerLinearMirror; + if (fill->spread() == FillSpread::Repeat) sampler = context.samplerLinearRepeat; + // update bind group + context.pipelines->layouts.releaseBindGroup(bindGroupGradient); + bindGroupGradient = context.pipelines->layouts.createBindGroupTexSampledBuff1Un( + sampler, texViewGradient, bufferGroupGradient); + } skip = false; } else if ((flags & (RenderUpdateFlag::Color)) && !fill) { rasterType = WgRenderRasterType::Solid; @@ -245,12 +253,11 @@ void WgRenderSettings::update(WgContext& context, const Fill* fill, const uint8_ void WgRenderSettings::release(WgContext& context) { context.pipelines->layouts.releaseBindGroup(bindGroupSolid); - context.pipelines->layouts.releaseBindGroup(bindGroupLinear); - context.pipelines->layouts.releaseBindGroup(bindGroupRadial); + context.pipelines->layouts.releaseBindGroup(bindGroupGradient); context.releaseBuffer(bufferGroupSolid); - context.releaseBuffer(bufferGroupLinear); - context.releaseBuffer(bufferGroupRadial); - + context.releaseBuffer(bufferGroupGradient); + context.releaseTexture(texGradient); + context.releaseTextureView(texViewGradient); }; //*********************************************************************** diff --git a/src/renderer/wg_engine/tvgWgRenderData.h b/src/renderer/wg_engine/tvgWgRenderData.h index 48ba0b87..f0837be9 100755 --- a/src/renderer/wg_engine/tvgWgRenderData.h +++ b/src/renderer/wg_engine/tvgWgRenderData.h @@ -78,11 +78,11 @@ enum class WgRenderRasterType { Solid = 0, Gradient, Image }; struct WgRenderSettings { WGPUBuffer bufferGroupSolid{}; - WGPUBuffer bufferGroupLinear{}; - WGPUBuffer bufferGroupRadial{}; WGPUBindGroup bindGroupSolid{}; - WGPUBindGroup bindGroupLinear{}; - WGPUBindGroup bindGroupRadial{}; + WGPUTexture texGradient{}; + WGPUTextureView texViewGradient{}; + WGPUBuffer bufferGroupGradient{}; + WGPUBindGroup bindGroupGradient{}; WgRenderSettingsType fillType{}; WgRenderRasterType rasterType{}; bool skip{}; diff --git a/src/renderer/wg_engine/tvgWgRenderTarget.cpp b/src/renderer/wg_engine/tvgWgRenderTarget.cpp index 0e58dd43..14a1710f 100755 --- a/src/renderer/wg_engine/tvgWgRenderTarget.cpp +++ b/src/renderer/wg_engine/tvgWgRenderTarget.cpp @@ -30,7 +30,7 @@ void WgRenderStorage::initialize(WgContext& context, uint32_t width, uint32_t he texView = context.createTextureView(texture); bindGroupRead = context.pipelines->layouts.createBindGroupStrorage1RO(texView); bindGroupWrite = context.pipelines->layouts.createBindGroupStrorage1WO(texView); - bindGroupTexure = context.pipelines->layouts.createBindGroupTexSampled(context.samplerNearest, texView); + bindGroupTexure = context.pipelines->layouts.createBindGroupTexSampled(context.samplerNearestRepeat, texView); } diff --git a/src/renderer/wg_engine/tvgWgRenderer.cpp b/src/renderer/wg_engine/tvgWgRenderer.cpp index 9ca23b88..b14910a9 100755 --- a/src/renderer/wg_engine/tvgWgRenderer.cpp +++ b/src/renderer/wg_engine/tvgWgRenderer.cpp @@ -21,7 +21,6 @@ */ #include "tvgWgRenderer.h" -#include WgRenderer::WgRenderer() { @@ -135,7 +134,7 @@ RenderData WgRenderer::prepare(Surface* surface, const RenderMesh* mesh, RenderD renderDataPicture->meshData.update(mContext, &geometryData); renderDataPicture->imageData.update(mContext, surface); renderDataPicture->bindGroupPicture = mContext.pipelines->layouts.createBindGroupTexSampled( - mContext.samplerLinear, renderDataPicture->imageData.textureView + mContext.samplerLinearRepeat, renderDataPicture->imageData.textureView ); } diff --git a/src/renderer/wg_engine/tvgWgShaderSrc.cpp b/src/renderer/wg_engine/tvgWgShaderSrc.cpp index 5407ef9f..5e2aaa5c 100755 --- a/src/renderer/wg_engine/tvgWgShaderSrc.cpp +++ b/src/renderer/wg_engine/tvgWgShaderSrc.cpp @@ -29,24 +29,15 @@ // graphics shader source: stencil //************************************************************************ -const char* cShaderSrc_Stencil = WG_SHADER_SOURCE( -// vertex input -struct VertexInput { - @location(0) position: vec2f -}; +const char* cShaderSrc_Stencil = R"( +struct VertexInput { @location(0) position: vec2f }; +struct VertexOutput { @builtin(position) position: vec4f }; -// vertex output -struct VertexOutput { - @builtin(position) position: vec4f -}; - -// uniforms @group(0) @binding(0) var uViewMat : mat4x4f; @group(1) @binding(0) var uModelMat : mat4x4f; @vertex fn vs_main(in: VertexInput) -> VertexOutput { - // fill output var out: VertexOutput; out.position = uViewMat * uModelMat * vec4f(in.position.xy, 0.0, 1.0); return out; @@ -56,40 +47,23 @@ fn vs_main(in: VertexInput) -> VertexOutput { fn fs_main(in: VertexOutput) -> @location(0) vec4f { return vec4f(0.0, 0.0, 0.0, 1.0); } -); +)"; //************************************************************************ // graphics shader source: solid //************************************************************************ -const char* cShaderSrc_Solid = WG_SHADER_SOURCE( -// vertex input -struct VertexInput { - @location(0) position: vec2f -}; +const char* cShaderSrc_Solid = R"( +struct VertexInput { @location(0) position: vec2f }; +struct VertexOutput { @builtin(position) position: vec4f }; -// BlendSettings -struct BlendSettings { - format : u32, // ColorSpace - dummy0 : f32, - dummy1 : f32, - opacity : f32 -}; - -// vertex output -struct VertexOutput { - @builtin(position) position: vec4f -}; - -// uniforms -@group(0) @binding(0) var uViewMat : mat4x4f; +@group(0) @binding(0) var uViewMat : mat4x4f; @group(1) @binding(0) var uModelMat : mat4x4f; -@group(1) @binding(1) var uBlendSettings : BlendSettings; -@group(2) @binding(0) var uSolidColor : vec4f; +@group(1) @binding(1) var uBlendSettings : vec4f; +@group(2) @binding(0) var uSolidColor : vec4f; @vertex fn vs_main(in: VertexInput) -> VertexOutput { - // fill output var out: VertexOutput; out.position = uViewMat * uModelMat * vec4f(in.position.xy, 0.0, 1.0); return out; @@ -97,64 +71,31 @@ fn vs_main(in: VertexInput) -> VertexOutput { @fragment fn fs_main(in: VertexOutput) -> @location(0) vec4f { - // resulting color - var color = vec4(1.0); - - // get color - color = uSolidColor; - - let alpha: f32 = color.a * uBlendSettings.opacity; - return vec4f(color.rgb * alpha, alpha); + let Sc = uSolidColor.rgb; + let Sa = uSolidColor.a; + let So = uBlendSettings.a; + return vec4f(Sc * Sa * So, Sa * So); } -); +)"; //************************************************************************ // graphics shader source: linear //************************************************************************ -const char* cShaderSrc_Linear = WG_SHADER_SOURCE( -// vertex input -struct VertexInput { - @location(0) position: vec2f -}; - -// BlendSettings -struct BlendSettings { - format : u32, // ColorSpace - dummy0 : f32, - dummy1 : f32, - opacity : f32 -}; - -// LinearGradient -const MAX_LINEAR_GRADIENT_STOPS = 32; -const MAX_LINEAR_GRADIENT_STOPS_PACKED = MAX_LINEAR_GRADIENT_STOPS / 4; -struct LinearGradient { - nStops : u32, - spread : u32, - dummy0 : u32, - dummy1 : u32, - gradStartPos : vec2f, - gradEndPos : vec2f, - stopPoints : array, - stopColors : array -}; - -// vertex output -struct VertexOutput { - @builtin(position) position : vec4f, - @location(0) vScreenCoord : vec2f -}; +const char* cShaderSrc_Linear = R"( +struct VertexInput { @location(0) position: vec2f }; +struct VertexOutput { @builtin(position) position : vec4f, @location(0) vScreenCoord : vec2f }; // uniforms -@group(0) @binding(0) var uViewMat : mat4x4f; -@group(1) @binding(0) var uModelMat : mat4x4f; -@group(1) @binding(1) var uBlendSettings : BlendSettings; -@group(2) @binding(0) var uLinearGradient : LinearGradient; +@group(0) @binding(0) var uViewMat : mat4x4f; +@group(1) @binding(0) var uModelMat : mat4x4f; +@group(1) @binding(1) var uBlendSettings : vec4f; +@group(2) @binding(0) var uSamplerGrad : sampler; +@group(2) @binding(1) var uTextureGrad : texture_2d; +@group(2) @binding(2) var uSettingGrad : vec4f; @vertex fn vs_main(in: VertexInput) -> VertexOutput { - // fill output var out: VertexOutput; out.position = uViewMat * uModelMat * vec4f(in.position.xy, 0.0, 1.0); out.vScreenCoord = in.position.xy; @@ -163,101 +104,34 @@ fn vs_main(in: VertexInput) -> VertexOutput { @fragment fn fs_main(in: VertexOutput) -> @location(0) vec4f { - // resulting color - var color = vec4(1.0); - - let pos: vec2f = in.vScreenCoord; - let st: vec2f = uLinearGradient.gradStartPos; - let ed: vec2f = uLinearGradient.gradEndPos; - - let ba: vec2f = ed - st; - - // get interpolation factor - var t: f32 = dot(pos - st, ba) / dot(ba, ba); - - // fill spread - switch uLinearGradient.spread { - /* Pad */ case 0u: { t = clamp(t, 0.0, 1.0); } - /* Reflect */ case 1u: { t = select(1.0 - fract(t), fract(t), u32(t) % 2 == 0); } - /* Repeat */ case 2u: { t = fract(t); } - default: { t = t; } - } - - // get stops count - let last: i32 = i32(uLinearGradient.nStops) - 1; - - // closer than first stop - if (t <= uLinearGradient.stopPoints[0][0]) { - color = uLinearGradient.stopColors[0]; - } - - // further than last stop - if (t >= uLinearGradient.stopPoints[last/4][last%4]) { - color = uLinearGradient.stopColors[last]; - } - - // look in the middle - for (var i = 0i; i < last; i++) { - let strt = uLinearGradient.stopPoints[i/4][i%4]; - let stop = uLinearGradient.stopPoints[(i+1)/4][(i+1)%4]; - if ((t >= strt) && (t < stop)) { - let step: f32 = (t - strt) / (stop - strt); - color = mix(uLinearGradient.stopColors[i], uLinearGradient.stopColors[i+1], step); - } - } - - let alpha: f32 = color.a * uBlendSettings.opacity; - return vec4f(color.rgb * alpha, alpha); + let pos = in.vScreenCoord; + let st = uSettingGrad.xy; + let ed = uSettingGrad.zw; + let ba = ed - st; + let t = dot(pos - st, ba) / dot(ba, ba); + let Sc = textureSample(uTextureGrad, uSamplerGrad, vec2f(t, 0.5)); + let So = uBlendSettings.a; + return vec4f(Sc.rgb * Sc.a * So, Sc.a * So); } -); +)"; //************************************************************************ // graphics shader source: radial //************************************************************************ -const char* cShaderSrc_Radial = WG_SHADER_SOURCE( -// vertex input -struct VertexInput { - @location(0) position: vec2f -}; +const char* cShaderSrc_Radial = R"( +struct VertexInput { @location(0) position: vec2f }; +struct VertexOutput { @builtin(position) position : vec4f, @location(0) vScreenCoord : vec2f }; -// BlendSettings -struct BlendSettings { - format : u32, // ColorSpace - dummy0 : f32, - dummy1 : f32, - opacity : f32 -}; - -// RadialGradient -const MAX_RADIAL_GRADIENT_STOPS = 32; -const MAX_RADIAL_GRADIENT_STOPS_PACKED = MAX_RADIAL_GRADIENT_STOPS / 4; -struct RadialGradient { - nStops : u32, - spread : u32, - dummy0 : u32, - dummy1 : u32, - centerPos : vec2f, - radius : vec2f, - stopPoints : array, - stopColors : array -}; - -// vertex output -struct VertexOutput { - @builtin(position) position : vec4f, - @location(0) vScreenCoord : vec2f -}; - -// uniforms -@group(0) @binding(0) var uViewMat : mat4x4f; -@group(1) @binding(0) var uModelMat : mat4x4f; -@group(1) @binding(1) var uBlendSettings : BlendSettings; -@group(2) @binding(0) var uRadialGradient : RadialGradient; +@group(0) @binding(0) var uViewMat : mat4x4f; +@group(1) @binding(0) var uModelMat : mat4x4f; +@group(1) @binding(1) var uBlendSettings : vec4f; +@group(2) @binding(0) var uSamplerGrad : sampler; +@group(2) @binding(1) var uTextureGrad : texture_2d; +@group(2) @binding(2) var uSettingGrad : vec4f; @vertex fn vs_main(in: VertexInput) -> VertexOutput { - // fill output var out: VertexOutput; out.position = uViewMat * uModelMat * vec4f(in.position.xy, 0.0, 1.0); out.vScreenCoord = in.position.xy; @@ -266,82 +140,29 @@ fn vs_main(in: VertexInput) -> VertexOutput { @fragment fn fs_main(in: VertexOutput) -> @location(0) vec4f { - // resulting color - var color = vec4(1.0); - - // get interpolation factor - var t: f32 = distance(uRadialGradient.centerPos, in.vScreenCoord) / uRadialGradient.radius.x; - - // fill spread - switch uRadialGradient.spread { - /* Pad */ case 0u: { t = clamp(t, 0.0, 1.0); } - /* Reflect */ case 1u: { t = select(1.0 - fract(t), fract(t), u32(t) % 2 == 0); } - /* Repeat */ case 2u: { t = fract(t); } - default: { t = t; } - } - - // get stops count - let last: i32 = i32(uRadialGradient.nStops) - 1; - - // closer than first stop - if (t <= uRadialGradient.stopPoints[0][0]) { - color = uRadialGradient.stopColors[0]; - } - - // further than last stop - if (t >= uRadialGradient.stopPoints[last/4][last%4]) { - color = uRadialGradient.stopColors[last]; - } - - // look in the middle - for (var i = 0i; i < last; i++) { - let strt = uRadialGradient.stopPoints[i/4][i%4]; - let stop = uRadialGradient.stopPoints[(i+1)/4][(i+1)%4]; - if ((t >= strt) && (t < stop)) { - let step: f32 = (t - strt) / (stop - strt); - color = mix(uRadialGradient.stopColors[i], uRadialGradient.stopColors[i+1], step); - } - } - - let alpha: f32 = color.a * uBlendSettings.opacity; - return vec4f(color.rgb * alpha, alpha); + let t: f32 = distance(uSettingGrad.zw, in.vScreenCoord) / uSettingGrad.r; + let Sc = textureSample(uTextureGrad, uSamplerGrad, vec2f(t, 0.5)); + let So = uBlendSettings.a; + return vec4f(Sc.rgb * Sc.a * So, Sc.a * So); } -); +)"; //************************************************************************ // graphics shader source: image //************************************************************************ -const char* cShaderSrc_Image = WG_SHADER_SOURCE( -// vertex input -struct VertexInput { - @location(0) position: vec2f, - @location(1) texCoord: vec2f -}; +const char* cShaderSrc_Image = R"( +struct VertexInput { @location(0) position: vec2f, @location(1) texCoord: vec2f }; +struct VertexOutput { @builtin(position) position: vec4f, @location(0) texCoord: vec2f }; -// BlendSettings -struct BlendSettings { - format : u32, // ColorSpace - dummy0 : f32, - dummy1 : f32, - opacity : f32 -}; - -// vertex output -struct VertexOutput { - @builtin(position) position: vec4f, - @location(0) texCoord: vec2f -}; - -@group(0) @binding(0) var uViewMat : mat4x4f; +@group(0) @binding(0) var uViewMat : mat4x4f; @group(1) @binding(0) var uModelMat : mat4x4f; -@group(1) @binding(1) var uBlendSettings : BlendSettings; -@group(2) @binding(0) var uSampler : sampler; -@group(2) @binding(1) var uTextureView : texture_2d; +@group(1) @binding(1) var uBlendSettings : vec4f; +@group(2) @binding(0) var uSampler : sampler; +@group(2) @binding(1) var uTextureView : texture_2d; @vertex fn vs_main(in: VertexInput) -> VertexOutput { - // fill output var out: VertexOutput; out.position = uViewMat * uModelMat * vec4f(in.position.xy, 0.0, 1.0); out.texCoord = in.texCoord; @@ -350,26 +171,24 @@ fn vs_main(in: VertexInput) -> VertexOutput { @fragment fn fs_main(in: VertexOutput) -> @location(0) vec4f { - var color: vec4f = textureSample(uTextureView, uSampler, in.texCoord.xy); - var result: vec4f = color; - var format: u32 = uBlendSettings.format; - if (format == 1u) { /* FMT_ARGB8888 */ - result = color.bgra; - } else if (format == 2u) { /* FMT_ABGR8888S */ - result = color.rgba; - } else if (format == 3u) { /* FMT_ARGB8888S */ - result = color.bgra; + var Sc: vec4f = textureSample(uTextureView, uSampler, in.texCoord.xy); + let So: f32 = uBlendSettings.a; + switch u32(uBlendSettings.r) { + case 0u: { Sc = Sc.rgba; } + case 1u: { Sc = Sc.bgra; } + case 2u: { Sc = Sc.rgba; } + case 3u: { Sc = Sc.bgra; } + default: {} } - return result * uBlendSettings.opacity; + return Sc * So; }; -); +)"; //************************************************************************ -// compute shaders blend +// compute shader source: merge clip path masks //************************************************************************ - -const char* cShaderSrc_MergeMasks = WG_SHADER_SOURCE( +const char* cShaderSrc_MergeMasks = R"( @group(0) @binding(0) var imageMsk0 : texture_storage_2d; @group(1) @binding(0) var imageMsk1 : texture_storage_2d; @group(2) @binding(0) var imageTrg : texture_storage_2d; @@ -380,326 +199,353 @@ fn cs_main(@builtin(global_invocation_id) id: vec3u) { let colorMsk1 = textureLoad(imageMsk1, id.xy); textureStore(imageTrg, id.xy, colorMsk0 * colorMsk1); } -); - +)"; //************************************************************************ -// compute shaders blend +// compute shader source: blend //************************************************************************ -std::string cBlendHeader = WG_SHADER_SOURCE( +const char* cShaderSrc_BlendHeader_Solid = R"( @group(0) @binding(0) var imageSrc : texture_storage_2d; @group(1) @binding(0) var imageDst : texture_storage_2d; @group(2) @binding(0) var imageTgt : texture_storage_2d; @group(3) @binding(0) var So : f32; -@compute @workgroup_size(8, 8) -fn cs_main(@builtin(global_invocation_id) id: vec3u) { +struct FragData { Sc: vec3f, Sa: f32, Dc: vec3f, Da: f32, skip: bool }; +fn getFragData(id: vec2u) -> FragData { + var data: FragData; + data.skip = true; let colorSrc = textureLoad(imageSrc, id.xy); - if (colorSrc.a == 0.0) { return; } + if (colorSrc.a == 0.0) { return data; } let colorDst = textureLoad(imageDst, id.xy); + data.Sc = colorSrc.rgb; + data.Sa = colorSrc.a; + data.Dc = colorDst.rgb; + data.Da = colorDst.a; + data.skip = false; + return data; +}; - 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; -); +fn postProcess(d: FragData, R: vec4f) -> vec4f { return R; }; +)"; -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 char* cShaderSrc_BlendHeader_Gradient = R"( +@group(0) @binding(0) var imageSrc : texture_storage_2d; +@group(1) @binding(0) var imageDst : texture_storage_2d; +@group(2) @binding(0) var imageTgt : texture_storage_2d; +@group(3) @binding(0) var So : f32; -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); -); +struct FragData { Sc: vec3f, Sa: f32, Dc: vec3f, Da: f32, skip: bool }; +fn getFragData(id: vec2u) -> FragData { + var data: FragData; + data.skip = true; + let colorSrc = textureLoad(imageSrc, id.xy); + if (colorSrc.a == 0.0) { return data; } + let colorDst = textureLoad(imageDst, id.xy); + data.Sc = colorSrc.rgb; + data.Sa = colorSrc.a; + data.Dc = colorDst.rgb; + data.Da = colorDst.a; + data.skip = false; + data.Sc = data.Sc + data.Dc * (1.0 - data.Sa); + data.Sa = data.Sa + data.Da * (1.0 - data.Sa); + return data; +}; -const std::string cBlendFooter = WG_SHADER_SOURCE( +fn postProcess(d: FragData, R: vec4f) -> vec4f { return R; }; +)"; + +const char* cShaderSrc_BlendHeader_Image = R"( +@group(0) @binding(0) var imageSrc : texture_storage_2d; +@group(1) @binding(0) var imageDst : texture_storage_2d; +@group(2) @binding(0) var imageTgt : texture_storage_2d; +@group(3) @binding(0) var So : f32; + +struct FragData { Sc: vec3f, Sa: f32, Dc: vec3f, Da: f32, skip: bool }; +fn getFragData(id: vec2u) -> FragData { + var data: FragData; + data.skip = true; + let colorSrc = textureLoad(imageSrc, id.xy); + if (colorSrc.a == 0.0) { return data; } + let colorDst = textureLoad(imageDst, id.xy); + data.Sc = colorSrc.rgb; + data.Sa = colorSrc.a; + data.Dc = colorDst.rgb; + data.Da = colorDst.a; + data.skip = false; + data.Sc = data.Sc * So; + data.Sa = data.Sa * So; + return data; +}; + +fn postProcess(d: FragData, R: vec4f) -> vec4f { + return mix(vec4(d.Dc, d.Da), R, d.Sa); +}; +)"; + +const char* cShaderSrc_Blend_Funcs = R"( +@compute @workgroup_size(8, 8) +fn cs_main_Normal(@builtin(global_invocation_id) id: vec3u) { + let d: FragData = getFragData(id.xy); + if (d.skip) { return; } + let Rc = d.Sc + d.Dc * (1.0 - d.Sa); + let Ra = d.Sa + d.Da * (1.0 - d.Sa); 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() +@compute @workgroup_size(8, 8) +fn cs_main_Add(@builtin(global_invocation_id) id: vec3u) { + let d: FragData = getFragData(id.xy); + if (d.skip) { return; } + let Rc = d.Sc + d.Dc; + textureStore(imageTgt, id.xy, postProcess(d, vec4f(Rc, 1.0))); }; -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 @workgroup_size(8, 8) +fn cs_main_Screen(@builtin(global_invocation_id) id: vec3u) { + let d: FragData = getFragData(id.xy); + if (d.skip) { return; } + let Rc = d.Sc + d.Dc - d.Sc * d.Dc; + textureStore(imageTgt, id.xy, postProcess(d, vec4f(Rc, 1.0))); }; +@compute @workgroup_size(8, 8) +fn cs_main_Multiply(@builtin(global_invocation_id) id: vec3u) { + let d: FragData = getFragData(id.xy); + if (d.skip) { return; } + let Rc = d.Sc * d.Dc; + textureStore(imageTgt, id.xy, postProcess(d, vec4f(Rc, 1.0))); +}; + +@compute @workgroup_size(8, 8) +fn cs_main_Overlay(@builtin(global_invocation_id) id: vec3u) { + let d: FragData = getFragData(id.xy); + if (d.skip) { return; } + var Rc: vec3f; + Rc.r = select(1.0 - min(1.0, 2 * (1 - d.Sc.r) * (1 - d.Dc.r)), min(1.0, 2 * d.Sc.r * d.Dc.r), (d.Dc.r < 0.5)); + Rc.g = select(1.0 - min(1.0, 2 * (1 - d.Sc.g) * (1 - d.Dc.g)), min(1.0, 2 * d.Sc.g * d.Dc.g), (d.Dc.g < 0.5)); + Rc.b = select(1.0 - min(1.0, 2 * (1 - d.Sc.b) * (1 - d.Dc.b)), min(1.0, 2 * d.Sc.b * d.Dc.b), (d.Dc.b < 0.5)); + textureStore(imageTgt, id.xy, postProcess(d, vec4f(Rc, 1.0))); +}; + +@compute @workgroup_size(8, 8) +fn cs_main_Difference(@builtin(global_invocation_id) id: vec3u) { + let d: FragData = getFragData(id.xy); + if (d.skip) { return; } + let Rc = abs(d.Dc - d.Sc); + textureStore(imageTgt, id.xy, postProcess(d, vec4f(Rc, 1.0))); +}; + +@compute @workgroup_size(8, 8) +fn cs_main_Exclusion(@builtin(global_invocation_id) id: vec3u) { + let d: FragData = getFragData(id.xy); + if (d.skip) { return; } + let One = vec3f(1.0, 1.0, 1.0); + let Rc = min(One, d.Sc + d.Dc - min(One, 2 * d.Sc * d.Dc)); + textureStore(imageTgt, id.xy, postProcess(d, vec4f(Rc, 1.0))); +}; + +@compute @workgroup_size(8, 8) +fn cs_main_SrcOver(@builtin(global_invocation_id) id: vec3u) { + let d: FragData = getFragData(id.xy); + if (d.skip) { return; } + textureStore(imageTgt, id.xy, postProcess(d, vec4f(d.Sc, d.Sa))); +}; + +@compute @workgroup_size(8, 8) +fn cs_main_Darken(@builtin(global_invocation_id) id: vec3u) { + let d: FragData = getFragData(id.xy); + if (d.skip) { return; } + let Rc = min(d.Sc, d.Dc); + textureStore(imageTgt, id.xy, postProcess(d, vec4f(Rc, 1.0))); +}; + +@compute @workgroup_size(8, 8) +fn cs_main_Lighten(@builtin(global_invocation_id) id: vec3u) { + let d: FragData = getFragData(id.xy); + if (d.skip) { return; } + let Rc = max(d.Sc, d.Dc); + textureStore(imageTgt, id.xy, postProcess(d, vec4f(Rc, 1.0))); +}; + +@compute @workgroup_size(8, 8) +fn cs_main_ColorDodge(@builtin(global_invocation_id) id: vec3u) { + let d: FragData = getFragData(id.xy); + if (d.skip) { return; } + var Rc: vec3f; + Rc.r = select(d.Dc.r, (d.Dc.r * 255.0 / (255.0 - d.Sc.r * 255.0))/255.0, (1.0 - d.Sc.r > 0.0)); + Rc.g = select(d.Dc.g, (d.Dc.g * 255.0 / (255.0 - d.Sc.g * 255.0))/255.0, (1.0 - d.Sc.g > 0.0)); + Rc.b = select(d.Dc.b, (d.Dc.b * 255.0 / (255.0 - d.Sc.b * 255.0))/255.0, (1.0 - d.Sc.b > 0.0)); + textureStore(imageTgt, id.xy, postProcess(d, vec4f(Rc, 1.0))); +}; + +@compute @workgroup_size(8, 8) +fn cs_main_ColorBurn(@builtin(global_invocation_id) id: vec3u) { + let d: FragData = getFragData(id.xy); + if (d.skip) { return; } + var Rc: vec3f; + Rc.r = select(1.0 - d.Dc.r, (255.0 - (255.0 - d.Dc.r * 255.0) / (d.Sc.r * 255.0)) / 255.0, (d.Sc.r > 0.0)); + Rc.g = select(1.0 - d.Dc.g, (255.0 - (255.0 - d.Dc.g * 255.0) / (d.Sc.g * 255.0)) / 255.0, (d.Sc.g > 0.0)); + Rc.b = select(1.0 - d.Dc.b, (255.0 - (255.0 - d.Dc.b * 255.0) / (d.Sc.b * 255.0)) / 255.0, (d.Sc.b > 0.0)); + textureStore(imageTgt, id.xy, postProcess(d, vec4f(Rc, 1.0))); +}; + +@compute @workgroup_size(8, 8) +fn cs_main_HardLight(@builtin(global_invocation_id) id: vec3u) { + let d: FragData = getFragData(id.xy); + if (d.skip) { return; } + var Rc: vec3f; + Rc.r = select(1.0 - min(1.0, 2 * (1 - d.Sc.r) * (1 - d.Dc.r)), min(1.0, 2 * d.Sc.r * d.Dc.r), (d.Sc.r < 0.5)); + Rc.g = select(1.0 - min(1.0, 2 * (1 - d.Sc.g) * (1 - d.Dc.g)), min(1.0, 2 * d.Sc.g * d.Dc.g), (d.Sc.g < 0.5)); + Rc.b = select(1.0 - min(1.0, 2 * (1 - d.Sc.b) * (1 - d.Dc.b)), min(1.0, 2 * d.Sc.b * d.Dc.b), (d.Sc.b < 0.5)); + textureStore(imageTgt, id.xy, postProcess(d, vec4f(Rc, 1.0))); +}; + +@compute @workgroup_size(8, 8) +fn cs_main_SoftLight(@builtin(global_invocation_id) id: vec3u) { + let d: FragData = getFragData(id.xy); + if (d.skip) { return; } + let One = vec3f(1.0, 1.0, 1.0); + let Rc = min(One, (One - 2 * d.Sc) * d.Dc * d.Dc + 2.0 * d.Sc * d.Dc); + textureStore(imageTgt, id.xy, postProcess(d, vec4f(Rc, 1.0))); +}; +)"; + //************************************************************************ -// compute shaders compose +// compute shader source: compose //************************************************************************ -std::string cComposeHeader = WG_SHADER_SOURCE( +const char* cShaderSrc_Compose = R"( @group(0) @binding(0) var imageSrc : texture_storage_2d; @group(1) @binding(0) var imageMsk : texture_storage_2d; @group(2) @binding(0) var imageDst : texture_storage_2d; @group(3) @binding(0) var imageTgt : texture_storage_2d; -@compute @workgroup_size(8, 8) -fn cs_main(@builtin(global_invocation_id) id: vec3u) { +struct FragData { Sc: vec3f, Sa: f32, Mc: vec3f, Ma: f32, Dc: vec3f, Da: f32 }; +fn getFragData(id: vec2u) -> FragData { let colorSrc = textureLoad(imageSrc, id.xy); let colorMsk = textureLoad(imageMsk, id.xy); let colorDst = textureLoad(imageDst, id.xy); - - var 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; -); - -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); -); - -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 cComposeEquation_LightenMask = WG_SHADER_SOURCE( - Rc = Sc; - Ra = select(Ma, Sa, Sa > Ma); -); -std::string cComposeEquation_DarkenMask = WG_SHADER_SOURCE( - Rc = Sc; - Ra = select(Sa, Ma, Sa > Ma); -); - -std::string cComposeFooter = WG_SHADER_SOURCE( - textureStore(imageTgt, id.xy, vec4f(Rc, Ra)); -} -); - -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; -std::string compose_LightenMask = cComposeHeader + cComposeEquation_LightenMask + cComposeFooter; -std::string compose_DarkenMask = cComposeHeader + cComposeEquation_DarkenMask + cComposeFooter; - -const char* cShaderSrc_Compose[12] { - 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(), - compose_LightenMask.c_str(), - compose_DarkenMask.c_str() + var data: FragData; + data.Sc = colorSrc.rgb; + data.Sa = colorSrc.a; + data.Mc = colorMsk.rgb; + data.Ma = colorMsk.a; + data.Dc = colorDst.rgb; + data.Da = colorDst.a; + return data; }; -const char* cShaderSrc_Copy = WG_SHADER_SOURCE( +@compute @workgroup_size(8, 8) +fn cs_main_None(@builtin(global_invocation_id) id: vec3u) { + let d: FragData = getFragData(id.xy); + let Rc = d.Dc; + let Ra = d.Da; + textureStore(imageTgt, id.xy, vec4f(Rc, Ra)); +}; + +@compute @workgroup_size(8, 8) +fn cs_main_ClipPath(@builtin(global_invocation_id) id: vec3u) { + let d: FragData = getFragData(id.xy); + let Rc = d.Sc * d.Ma + d.Dc * (1.0 - d.Sa * d.Ma); + let Ra = d.Sa * d.Ma + d.Da * (1.0 - d.Sa * d.Ma); + textureStore(imageTgt, id.xy, vec4f(Rc, Ra)); +}; + +@compute @workgroup_size(8, 8) +fn cs_main_AlphaMask(@builtin(global_invocation_id) id: vec3u) { + let d: FragData = getFragData(id.xy); + let Rc = d.Sc * d.Ma + d.Dc * (1.0 - d.Sa * d.Ma); + let Ra = d.Sa * d.Ma + d.Da * (1.0 - d.Sa * d.Ma); + textureStore(imageTgt, id.xy, vec4f(Rc, Ra)); +}; + +@compute @workgroup_size(8, 8) +fn cs_main_InvAlphaMask(@builtin(global_invocation_id) id: vec3u) { + let d: FragData = getFragData(id.xy); + let Rc = d.Sc * (1.0 - d.Ma) + d.Dc * (1.0 - d.Sa * (1.0 - d.Ma)); + let Ra = d.Sa * (1.0 - d.Ma) + d.Da * (1.0 - d.Sa * (1.0 - d.Ma)); + textureStore(imageTgt, id.xy, vec4f(Rc, Ra)); +}; + +@compute @workgroup_size(8, 8) +fn cs_main_LumaMask(@builtin(global_invocation_id) id: vec3u) { + let d: FragData = getFragData(id.xy); + let luma: f32 = dot(d.Mc, vec3f(0.2125, 0.7154, 0.0721)); + let Rc = d.Sc * luma + d.Dc * (1.0 - d.Sa * luma); + let Ra = d.Sa * luma + d.Da * (1.0 - d.Sa * luma); + textureStore(imageTgt, id.xy, vec4f(Rc, Ra)); +}; + +@compute @workgroup_size(8, 8) +fn cs_main_InvLumaMask(@builtin(global_invocation_id) id: vec3u) { + let d: FragData = getFragData(id.xy); + let luma: f32 = dot(d.Mc, vec3f(0.2125, 0.7154, 0.0721)); + let Rc = d.Sc * (1.0 - luma) + d.Dc * (1.0 - d.Sa * (1.0 - luma)); + let Ra = d.Sa * (1.0 - luma) + d.Da * (1.0 - d.Sa * (1.0 - luma)); + textureStore(imageTgt, id.xy, vec4f(Rc, Ra)); +}; + +@compute @workgroup_size(8, 8) +fn cs_main_AddMask(@builtin(global_invocation_id) id: vec3u) { + let d: FragData = getFragData(id.xy); + let Rc = d.Sc; + let Ra = d.Sa + d.Ma * (1.0 - d.Sa); + textureStore(imageTgt, id.xy, vec4f(Rc, Ra)); +}; + +@compute @workgroup_size(8, 8) +fn cs_main_SubtractMask(@builtin(global_invocation_id) id: vec3u) { + let d: FragData = getFragData(id.xy); + let Rc = d.Sc; + let Ra = d.Sa * (1.0 - d.Ma); + textureStore(imageTgt, id.xy, vec4f(Rc, Ra)); +}; + +@compute @workgroup_size(8, 8) +fn cs_main_IntersectMask(@builtin(global_invocation_id) id: vec3u) { + let d: FragData = getFragData(id.xy); + let Rc = d.Sc; + let Ra = d.Sa * d.Ma; + textureStore(imageTgt, id.xy, vec4f(Rc, Ra)); +}; + +@compute @workgroup_size(8, 8) +fn cs_main_DifferenceMask(@builtin(global_invocation_id) id: vec3u) { + let d: FragData = getFragData(id.xy); + let Rc = d.Sc; + let Ra = d.Sa * (1.0 - d.Ma) + d.Ma * (1.0 - d.Sa); + textureStore(imageTgt, id.xy, vec4f(Rc, Ra)); +}; + +@compute @workgroup_size(8, 8) +fn cs_main_LightenMask(@builtin(global_invocation_id) id: vec3u) { + let d: FragData = getFragData(id.xy); + let Rc = d.Sc; + let Ra = max(d.Sa, d.Ma); + textureStore(imageTgt, id.xy, vec4f(Rc, Ra)); +}; + +@compute @workgroup_size(8, 8) +fn cs_main_DarkenMask(@builtin(global_invocation_id) id: vec3u) { + let d: FragData = getFragData(id.xy); + let Rc = d.Sc; + let Ra = min(d.Sa, d.Ma); + textureStore(imageTgt, id.xy, vec4f(Rc, Ra)); +}; +)"; + +//************************************************************************ +// compute shader source: copy +//************************************************************************ + +const char* cShaderSrc_Copy = R"( @group(0) @binding(0) var imageSrc : texture_storage_2d; @group(1) @binding(0) var imageDst : texture_storage_2d; @compute @workgroup_size(8, 8) fn cs_main(@builtin(global_invocation_id) id: vec3u) { textureStore(imageDst, id.xy, textureLoad(imageSrc, id.xy)); -} -); \ No newline at end of file +}; +)"; diff --git a/src/renderer/wg_engine/tvgWgShaderSrc.h b/src/renderer/wg_engine/tvgWgShaderSrc.h index 13c269c7..85a1980d 100755 --- a/src/renderer/wg_engine/tvgWgShaderSrc.h +++ b/src/renderer/wg_engine/tvgWgShaderSrc.h @@ -32,11 +32,11 @@ extern const char* cShaderSrc_Image; // 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[12]; +extern const char* cShaderSrc_BlendHeader_Solid; +extern const char* cShaderSrc_BlendHeader_Gradient; +extern const char* cShaderSrc_BlendHeader_Image; +extern const char* cShaderSrc_Blend_Funcs; +extern const char* cShaderSrc_Compose; extern const char* cShaderSrc_Copy; #endif // _TVG_WG_SHEDER_SRC_H_ diff --git a/src/renderer/wg_engine/tvgWgShaderTypes.cpp b/src/renderer/wg_engine/tvgWgShaderTypes.cpp index cc109f4c..14db7369 100755 --- a/src/renderer/wg_engine/tvgWgShaderTypes.cpp +++ b/src/renderer/wg_engine/tvgWgShaderTypes.cpp @@ -21,6 +21,7 @@ */ #include "tvgWgShaderTypes.h" +#include /////////////////////////////////////////////////////////////////////////////// // shader types @@ -92,10 +93,8 @@ WgShaderTypeBlendSettings::WgShaderTypeBlendSettings(const ColorSpace colorSpace void WgShaderTypeBlendSettings::update(const ColorSpace colorSpace, uint8_t o) { - format = (uint32_t)colorSpace; - dummy0 = 0.0f; - dummy1 = 0.0f; - opacity = o / 255.0f; + settings[0] = (uint32_t)colorSpace; + settings[3] = o / 255.0f; } @@ -114,53 +113,59 @@ void WgShaderTypeSolidColor::update(const uint8_t* c) } -WgShaderTypeLinearGradient::WgShaderTypeLinearGradient(const LinearGradient* linearGradient) -{ - update(linearGradient); -} - - -void WgShaderTypeLinearGradient::update(const LinearGradient* linearGradient) +void WgShaderTypeGradient::update(const LinearGradient* linearGradient) { + // update gradient data const Fill::ColorStop* stops = nullptr; auto stopCnt = linearGradient->colorStops(&stops); - - nStops = stopCnt; - spread = uint32_t(linearGradient->spread()); - - for (uint32_t i = 0; i < stopCnt; ++i) { - stopPoints[i] = stops[i].offset; - stopColors[i * 4 + 0] = stops[i].r / 255.f; - stopColors[i * 4 + 1] = stops[i].g / 255.f; - stopColors[i * 4 + 2] = stops[i].b / 255.f; - stopColors[i * 4 + 3] = stops[i].a / 255.f; - } - - linearGradient->linear(&startPos[0], &startPos[1], &endPos[0], &endPos[1]); -} + updateTexData(stops, stopCnt); + // update base points + linearGradient->linear(&settings[0], &settings[1], &settings[2], &settings[3]); +}; -WgShaderTypeRadialGradient::WgShaderTypeRadialGradient(const RadialGradient* radialGradient) -{ - update(radialGradient); -} - - -void WgShaderTypeRadialGradient::update(const RadialGradient* radialGradient) +void WgShaderTypeGradient::update(const RadialGradient* radialGradient) { + // update gradient data const Fill::ColorStop* stops = nullptr; auto stopCnt = radialGradient->colorStops(&stops); + updateTexData(stops, stopCnt); + // update base points + radialGradient->radial(&settings[2], &settings[3], &settings[0]); +}; - nStops = stopCnt; - spread = uint32_t(radialGradient->spread()); - for (uint32_t i = 0; i < stopCnt; ++i) { - stopPoints[i] = stops[i].offset; - stopColors[i * 4 + 0] = stops[i].r / 255.f; - stopColors[i * 4 + 1] = stops[i].g / 255.f; - stopColors[i * 4 + 2] = stops[i].b / 255.f; - stopColors[i * 4 + 3] = stops[i].a / 255.f; +void WgShaderTypeGradient::updateTexData(const Fill::ColorStop* stops, uint32_t stopCnt) +{ + assert(stops); + // head + uint32_t range_s = 0; + uint32_t range_e = uint32_t(stops[0].offset * (WG_TEXTURE_GRADIENT_SIZE-1)); + for (uint32_t ti = range_s; (ti < range_e) && (ti < WG_TEXTURE_GRADIENT_SIZE); ti++) { + texData[ti * 4 + 0] = stops[0].r; + texData[ti * 4 + 1] = stops[0].g; + texData[ti * 4 + 2] = stops[0].b; + texData[ti * 4 + 3] = stops[0].a; + } + // body + for (uint32_t di = 1; di < stopCnt; di++) { + range_s = uint32_t(stops[di-1].offset * (WG_TEXTURE_GRADIENT_SIZE-1)); + range_e = uint32_t(stops[di-0].offset * (WG_TEXTURE_GRADIENT_SIZE-1)); + for (uint32_t ti = range_s; (ti < range_e) && (ti < WG_TEXTURE_GRADIENT_SIZE); ti++) { + float t = float(ti - range_s) / (range_e - range_s); + texData[ti * 4 + 0] = uint8_t((1.0f - t) * stops[di-1].r + t * stops[di].r); + texData[ti * 4 + 1] = uint8_t((1.0f - t) * stops[di-1].g + t * stops[di].g); + texData[ti * 4 + 2] = uint8_t((1.0f - t) * stops[di-1].b + t * stops[di].b); + texData[ti * 4 + 3] = uint8_t((1.0f - t) * stops[di-1].a + t * stops[di].a); + } + } + // tail + range_s = uint32_t(stops[stopCnt-1].offset * (WG_TEXTURE_GRADIENT_SIZE-1)); + range_e = WG_TEXTURE_GRADIENT_SIZE; + for (uint32_t ti = range_s; ti < range_e; ti++) { + texData[ti * 4 + 0] = stops[stopCnt-1].r; + texData[ti * 4 + 1] = stops[stopCnt-1].g; + texData[ti * 4 + 2] = stops[stopCnt-1].b; + texData[ti * 4 + 3] = stops[stopCnt-1].a; } - - radialGradient->radial(¢erPos[0], ¢erPos[1], &radius[0]); } diff --git a/src/renderer/wg_engine/tvgWgShaderTypes.h b/src/renderer/wg_engine/tvgWgShaderTypes.h index f9cc1015..eddfc056 100755 --- a/src/renderer/wg_engine/tvgWgShaderTypes.h +++ b/src/renderer/wg_engine/tvgWgShaderTypes.h @@ -42,27 +42,17 @@ struct WgShaderTypeMat4x4f void update(size_t w, size_t h); }; -// struct BlendSettings { -// format : u32, // ColorSpace -// dummy0 : f32, -// dummy1 : f32, -// opacity : f32 -// }; +// vec4f struct WgShaderTypeBlendSettings { - uint32_t format{}; // ColorSpace - float dummy0{}; - float dummy1{}; - float opacity{}; + float settings[4]{}; WgShaderTypeBlendSettings() {}; WgShaderTypeBlendSettings(const ColorSpace colorSpace, uint8_t o); void update(const ColorSpace colorSpace, uint8_t o); }; -// struct SolidColor { -// color: vec4f -// }; +// vec4f struct WgShaderTypeSolidColor { float color[4]{}; @@ -71,58 +61,16 @@ struct WgShaderTypeSolidColor void update(const uint8_t* c); }; -// const MAX_LINEAR_GRADIENT_STOPS = 4; -// struct LinearGradient { -// nStops : u32, -// spread : u32, -// dummy0 : u32, -// dummy1 : u32, -// gradStartPos : vec2f, -// gradEndPos : vec2f, -// stopPoints : vec4f, -// stopColors : array -// }; -#define MAX_LINEAR_GRADIENT_STOPS 32 -struct WgShaderTypeLinearGradient +// sampler, texture, vec4f +#define WG_TEXTURE_GRADIENT_SIZE 512 +struct WgShaderTypeGradient { - uint32_t nStops{}; - uint32_t spread{}; - uint32_t dummy0{}; // align with WGSL struct - uint32_t dummy1{}; // align with WGSL struct - float startPos[2]{}; - float endPos[2]{}; - float stopPoints[MAX_LINEAR_GRADIENT_STOPS]{}; - float stopColors[4 * MAX_LINEAR_GRADIENT_STOPS]{}; + float settings[4]{}; + uint8_t texData[WG_TEXTURE_GRADIENT_SIZE * 4]; - WgShaderTypeLinearGradient(const LinearGradient* linearGradient); void update(const LinearGradient* linearGradient); -}; - -// const MAX_RADIAL_GRADIENT_STOPS = 4; -// struct RadialGradient { -// nStops : u32, -// spread : u32, -// dummy0 : u32, -// dummy1 : u32, -// centerPos : vec2f, -// radius : vec2f, -// stopPoints : vec4f, -// stopColors : array -// }; -#define MAX_RADIAL_GRADIENT_STOPS 32 -struct WgShaderTypeRadialGradient -{ - uint32_t nStops{}; - uint32_t spread{}; - uint32_t dummy0{}; // align with WGSL struct - uint32_t dummy1{}; // align with WGSL struct - float centerPos[2]{}; - float radius[2]{}; - float stopPoints[MAX_RADIAL_GRADIENT_STOPS]{}; - float stopColors[4 * MAX_RADIAL_GRADIENT_STOPS]{}; - - WgShaderTypeRadialGradient(const RadialGradient* radialGradient); void update(const RadialGradient* radialGradient); + void updateTexData(const Fill::ColorStop* stops, uint32_t stopCnt); }; #endif // _TVG_WG_SHADER_TYPES_H_