From 6c248f61743d2eee1d49db011b99b490927d1d2c Mon Sep 17 00:00:00 2001 From: The Roofer Dev Date: Fri, 26 Jun 2026 18:25:48 -0400 Subject: [PATCH] =?UTF-8?q?v1.1.6=20=E2=80=94=20NIS=20(NVIDIA=20Image=20Sc?= =?UTF-8?q?aling)=20scaling=20filter=20+=20HDR=20highlight=20whitening=20s?= =?UTF-8?q?lider?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- .../Configuration/ScalingFilter.cs | 1 + src/Ryujinx.Graphics.GAL/IWindow.cs | 2 +- .../Multithreading/ThreadedWindow.cs | 2 +- src/Ryujinx.Graphics.OpenGL/Window.cs | 2 +- .../Effects/AreaScalingFilter.cs | 5 +- .../Effects/FsrScalingFilter.cs | 10 +- .../Effects/IScalingFilter.cs | 4 +- .../Effects/NisScalingFilter.cs | 231 ++++ .../Effects/Shaders/FsrSharpening.glsl | 11 +- .../Effects/Shaders/FsrSharpeningHdr.spv | Bin 22388 -> 22952 bytes .../Effects/Shaders/NisScaler.h | 1023 +++++++++++++++++ .../Effects/Shaders/NisScaling.glsl | 88 ++ .../Effects/Shaders/NisScaling.spv | Bin 0 -> 56760 bytes .../Effects/Shaders/NisScalingHdr.glsl | 125 ++ .../Effects/Shaders/NisScalingHdr.spv | Bin 0 -> 59420 bytes .../Effects/Shaders/nis_coef.glsl | 157 +++ src/Ryujinx.Graphics.Vulkan/HelperShader.cs | 4 +- .../Ryujinx.Graphics.Vulkan.csproj | 2 + .../ColorBlitHdrFragmentShaderSource.frag | 13 +- .../SpirvBinaries/ColorBlitHdrFragment.spv | Bin 2952 -> 3520 bytes src/Ryujinx.Graphics.Vulkan/Window.cs | 21 +- src/Ryujinx.Graphics.Vulkan/WindowBase.cs | 2 +- src/Ryujinx/Systems/AppHost.cs | 8 +- .../Configuration/ConfigurationFileFormat.cs | 7 +- .../ConfigurationState.Migration.cs | 4 +- .../Configuration/ConfigurationState.Model.cs | 8 + .../Configuration/ConfigurationState.cs | 2 + .../UI/ViewModels/SettingsViewModel.cs | 15 + .../Views/Settings/SettingsGraphicsView.axaml | 24 + 29 files changed, 1749 insertions(+), 22 deletions(-) create mode 100644 src/Ryujinx.Graphics.Vulkan/Effects/NisScalingFilter.cs create mode 100644 src/Ryujinx.Graphics.Vulkan/Effects/Shaders/NisScaler.h create mode 100644 src/Ryujinx.Graphics.Vulkan/Effects/Shaders/NisScaling.glsl create mode 100644 src/Ryujinx.Graphics.Vulkan/Effects/Shaders/NisScaling.spv create mode 100644 src/Ryujinx.Graphics.Vulkan/Effects/Shaders/NisScalingHdr.glsl create mode 100644 src/Ryujinx.Graphics.Vulkan/Effects/Shaders/NisScalingHdr.spv create mode 100644 src/Ryujinx.Graphics.Vulkan/Effects/Shaders/nis_coef.glsl diff --git a/src/Ryujinx.Common/Configuration/ScalingFilter.cs b/src/Ryujinx.Common/Configuration/ScalingFilter.cs index 9040b1be0..c04a18a51 100644 --- a/src/Ryujinx.Common/Configuration/ScalingFilter.cs +++ b/src/Ryujinx.Common/Configuration/ScalingFilter.cs @@ -9,5 +9,6 @@ namespace Ryujinx.Common.Configuration Nearest, Fsr, Area, + Nis, } } diff --git a/src/Ryujinx.Graphics.GAL/IWindow.cs b/src/Ryujinx.Graphics.GAL/IWindow.cs index d5c048784..a364867d1 100644 --- a/src/Ryujinx.Graphics.GAL/IWindow.cs +++ b/src/Ryujinx.Graphics.GAL/IWindow.cs @@ -15,6 +15,6 @@ namespace Ryujinx.Graphics.GAL void SetScalingFilter(ScalingFilter type); void SetScalingFilterLevel(float level); void SetColorSpacePassthrough(bool colorSpacePassThroughEnabled); - void SetHdrMode(bool enabled, float paperWhite, float peak, float curve, float gamma, float blend); + void SetHdrMode(bool enabled, float paperWhite, float peak, float curve, float gamma, float blend, float whiten); } } diff --git a/src/Ryujinx.Graphics.GAL/Multithreading/ThreadedWindow.cs b/src/Ryujinx.Graphics.GAL/Multithreading/ThreadedWindow.cs index 220ab8a90..9d1d23a33 100644 --- a/src/Ryujinx.Graphics.GAL/Multithreading/ThreadedWindow.cs +++ b/src/Ryujinx.Graphics.GAL/Multithreading/ThreadedWindow.cs @@ -42,6 +42,6 @@ namespace Ryujinx.Graphics.GAL.Multithreading public void SetColorSpacePassthrough(bool colorSpacePassthroughEnabled) { } - public void SetHdrMode(bool enabled, float paperWhite, float peak, float curve, float gamma, float blend) { } + public void SetHdrMode(bool enabled, float paperWhite, float peak, float curve, float gamma, float blend, float whiten) { } } } diff --git a/src/Ryujinx.Graphics.OpenGL/Window.cs b/src/Ryujinx.Graphics.OpenGL/Window.cs index 2518c7d33..ab4e17c62 100644 --- a/src/Ryujinx.Graphics.OpenGL/Window.cs +++ b/src/Ryujinx.Graphics.OpenGL/Window.cs @@ -424,7 +424,7 @@ namespace Ryujinx.Graphics.OpenGL _updateScalingFilter = true; } - public void SetHdrMode(bool enabled, float paperWhite, float peak, float curve, float gamma, float blend) + public void SetHdrMode(bool enabled, float paperWhite, float peak, float curve, float gamma, float blend, float whiten) { // HDR output (scRGB inverse tone mapping) is only implemented on the Vulkan backend. } diff --git a/src/Ryujinx.Graphics.Vulkan/Effects/AreaScalingFilter.cs b/src/Ryujinx.Graphics.Vulkan/Effects/AreaScalingFilter.cs index fc6dc2a02..1e0f16686 100644 --- a/src/Ryujinx.Graphics.Vulkan/Effects/AreaScalingFilter.cs +++ b/src/Ryujinx.Graphics.Vulkan/Effects/AreaScalingFilter.cs @@ -55,6 +55,8 @@ namespace Ryujinx.Graphics.Vulkan.Effects ], scalingResourceLayout); } + public bool IsResolutionSupported(int srcWidth, int srcHeight, int dstWidth, int dstHeight) => true; + public void Run( TextureView view, CommandBufferScoped cbs, @@ -69,7 +71,8 @@ namespace Ryujinx.Graphics.Vulkan.Effects float peak = 12.5f, float curve = 4.0f, float gamma = 2.2f, - float blend = 0.5f) + float blend = 0.5f, + float whiten = 0.0f) { _pipeline.SetCommandBuffer(cbs); _pipeline.SetProgram(_scalingProgram); diff --git a/src/Ryujinx.Graphics.Vulkan/Effects/FsrScalingFilter.cs b/src/Ryujinx.Graphics.Vulkan/Effects/FsrScalingFilter.cs index 69d60e9d7..1fd14f1a2 100644 --- a/src/Ryujinx.Graphics.Vulkan/Effects/FsrScalingFilter.cs +++ b/src/Ryujinx.Graphics.Vulkan/Effects/FsrScalingFilter.cs @@ -87,6 +87,8 @@ namespace Ryujinx.Graphics.Vulkan.Effects ], sharpeningResourceLayout); } + public bool IsResolutionSupported(int srcWidth, int srcHeight, int dstWidth, int dstHeight) => true; + public void Run( TextureView view, CommandBufferScoped cbs, @@ -101,7 +103,8 @@ namespace Ryujinx.Graphics.Vulkan.Effects float peak = 12.5f, float curve = 4.0f, float gamma = 2.2f, - float blend = 0.5f) + float blend = 0.5f, + float whiten = 0.0f) { if (_intermediaryTexture == null || _intermediaryTexture.Info.Width != width @@ -159,14 +162,15 @@ namespace Ryujinx.Graphics.Vulkan.Effects // Non-HDR: [sharpening]. HDR: [sharpening, paperWhite, peak, curve, gamma, blend] (the HDR sharpening // shader reads the extra floats from the same uniform block to tone-map to scRGB). - int sharpeningCount = hdr ? 6 : 1; - Span sharpeningBufferData = stackalloc float[6]; + int sharpeningCount = hdr ? 7 : 1; + Span sharpeningBufferData = stackalloc float[7]; sharpeningBufferData[0] = 1.5f - (Level * 0.01f * 1.5f); sharpeningBufferData[1] = paperWhite; sharpeningBufferData[2] = peak; sharpeningBufferData[3] = curve; sharpeningBufferData[4] = gamma; sharpeningBufferData[5] = blend; + sharpeningBufferData[6] = whiten; using ScopedTemporaryBuffer sharpeningBuffer = _renderer.BufferManager.ReserveOrCreate(_renderer, cbs, sharpeningCount * sizeof(float)); sharpeningBuffer.Holder.SetDataUnchecked(sharpeningBuffer.Offset, sharpeningBufferData[..sharpeningCount]); diff --git a/src/Ryujinx.Graphics.Vulkan/Effects/IScalingFilter.cs b/src/Ryujinx.Graphics.Vulkan/Effects/IScalingFilter.cs index a12fb0a2e..c68ad7f45 100644 --- a/src/Ryujinx.Graphics.Vulkan/Effects/IScalingFilter.cs +++ b/src/Ryujinx.Graphics.Vulkan/Effects/IScalingFilter.cs @@ -7,6 +7,7 @@ namespace Ryujinx.Graphics.Vulkan.Effects internal interface IScalingFilter : IDisposable { float Level { get; set; } + bool IsResolutionSupported(int srcWidth, int srcHeight, int dstWidth, int dstHeight); void Run( TextureView view, CommandBufferScoped cbs, @@ -21,6 +22,7 @@ namespace Ryujinx.Graphics.Vulkan.Effects float peak = 12.5f, float curve = 4.0f, float gamma = 2.2f, - float blend = 0.5f); + float blend = 0.5f, + float whiten = 0.0f); } } diff --git a/src/Ryujinx.Graphics.Vulkan/Effects/NisScalingFilter.cs b/src/Ryujinx.Graphics.Vulkan/Effects/NisScalingFilter.cs new file mode 100644 index 000000000..21f7a58b2 --- /dev/null +++ b/src/Ryujinx.Graphics.Vulkan/Effects/NisScalingFilter.cs @@ -0,0 +1,231 @@ +using Ryujinx.Common; +using Ryujinx.Graphics.GAL; +using Ryujinx.Graphics.Shader; +using Ryujinx.Graphics.Shader.Translation; +using Silk.NET.Vulkan; +using System; +using Extent2D = Ryujinx.Graphics.GAL.Extents2D; +using Format = Silk.NET.Vulkan.Format; +using SamplerCreateInfo = Ryujinx.Graphics.GAL.SamplerCreateInfo; + +namespace Ryujinx.Graphics.Vulkan.Effects +{ + /// + /// NVIDIA Image Scaling (NIS) upscaler. Wraps the MIT-licensed NIS NVScaler compute + /// shader (see Effects/Shaders/NisScaler.h, nis_coef.glsl, NisScaling.glsl). + /// NVScaler is only defined for upscale ratios in [1x, 2x]; outside that the input + /// is simply stretched without the NIS quality benefit. SDR (rgba8) path for now; + /// an scRGB/HDR variant is a follow-up. + /// + internal class NisScalingFilter : IScalingFilter + { + // NIS shader block dimensions (must match NisScaling.glsl: NIS_BLOCK_WIDTH/HEIGHT). + private const int BlockWidth = 32; + private const int BlockHeight = 24; + + private const int ConfigFloats = 30; // NISConfig: 28 named values + 2 reserved. + + private readonly VulkanRenderer _renderer; + private PipelineHelperShader _pipeline; + private ISampler _sampler; + private ShaderCollection _scalingProgram; + private ShaderCollection _scalingProgramHdr; + private Device _device; + + // Slider value (0..100). Mapped to NIS sharpness (0..1) in Run. 50 = neutral. + public float Level { get; set; } = 50f; + + public NisScalingFilter(VulkanRenderer renderer, Device device) + { + _device = device; + _renderer = renderer; + + Initialize(); + } + + public void Dispose() + { + _pipeline.Dispose(); + _scalingProgram.Dispose(); + _scalingProgramHdr.Dispose(); + _sampler.Dispose(); + } + + public void Initialize() + { + _pipeline = new PipelineHelperShader(_renderer, _device); + _pipeline.Initialize(); + + byte[] scalingShader = EmbeddedResources.Read("Ryujinx.Graphics.Vulkan/Effects/Shaders/NisScaling.spv"); + + // Same layout as the other scaling filters: config UBO (2), input+sampler (1), output image (0/set 3). + ResourceLayout scalingResourceLayout = new ResourceLayoutBuilder() + .Add(ResourceStages.Compute, ResourceType.UniformBuffer, 2) + .Add(ResourceStages.Compute, ResourceType.TextureAndSampler, 1) + .Add(ResourceStages.Compute, ResourceType.Image, 0, true).Build(); + + _sampler = _renderer.CreateSampler(SamplerCreateInfo.Create(MinFilter.Linear, MagFilter.Linear)); + + _scalingProgram = _renderer.CreateProgramWithMinimalLayout([ + new ShaderSource(scalingShader, ShaderStage.Compute, TargetLanguage.Spirv) + ], scalingResourceLayout); + + byte[] scalingShaderHdr = EmbeddedResources.Read("Ryujinx.Graphics.Vulkan/Effects/Shaders/NisScalingHdr.spv"); + + // HDR variant adds the tone-map params UBO at binding 3 and outputs to rgba16f (scRGB). + ResourceLayout scalingResourceLayoutHdr = new ResourceLayoutBuilder() + .Add(ResourceStages.Compute, ResourceType.UniformBuffer, 2) + .Add(ResourceStages.Compute, ResourceType.UniformBuffer, 3) + .Add(ResourceStages.Compute, ResourceType.TextureAndSampler, 1) + .Add(ResourceStages.Compute, ResourceType.Image, 0, true).Build(); + + _scalingProgramHdr = _renderer.CreateProgramWithMinimalLayout([ + new ShaderSource(scalingShaderHdr, ShaderStage.Compute, TargetLanguage.Spirv) + ], scalingResourceLayoutHdr); + } + + // NIS NVScaler is only defined for upscale ratios in [1x, 2x] (kScale in [0.5, 1.0]). + // Outside that range - notably when the internal resolution is at/above the output, as + // with high-res mods or texture packs - its block tile cache reads out of bounds and + // corrupts the image. Report unsupported so Window.Present falls back to a normal blit. + public bool IsResolutionSupported(int srcWidth, int srcHeight, int dstWidth, int dstHeight) + { + float scaleX = (float)srcWidth / dstWidth; + float scaleY = (float)srcHeight / dstHeight; + return scaleX >= 0.5f && scaleX <= 1.0f && scaleY >= 0.5f && scaleY <= 1.0f; + } + + public void Run( + TextureView view, + CommandBufferScoped cbs, + Auto destinationTexture, + Format format, + int width, + int height, + Extent2D source, + Extent2D destination, + bool hdr = false, + float paperWhite = 2.5f, + float peak = 12.5f, + float curve = 4.0f, + float gamma = 2.2f, + float blend = 0.5f, + float whiten = 0.0f) + { + _pipeline.SetCommandBuffer(cbs); + _pipeline.SetProgram(hdr ? _scalingProgramHdr : _scalingProgram); + _pipeline.SetTextureAndSampler(ShaderStage.Compute, 1, view, _sampler); + + // v1: stretch the full source frame to the full output (swapchain) surface. + // Letterboxing/flip via source/destination extents is a follow-up. + uint inW = (uint)view.Width; + uint inH = (uint)view.Height; + uint outW = (uint)width; + uint outH = (uint)height; + + float sharpness = Math.Clamp(Level * 0.01f, 0f, 1f); + + Span cb = stackalloc float[ConfigFloats]; + BuildConfig(cb, sharpness, inW, inH, outW, outH); + + int rangeSize = ConfigFloats * sizeof(float); + using ScopedTemporaryBuffer buffer = _renderer.BufferManager.ReserveOrCreate(_renderer, cbs, rangeSize); + buffer.Holder.SetDataUnchecked(buffer.Offset, cb); + + // HDR tone-map params (paperWhite, peak, curve, gamma, blend, whiten), matching NisScalingHdr.glsl. + Span hdrData = stackalloc float[] { paperWhite, peak, curve, gamma, blend, whiten }; + using ScopedTemporaryBuffer hdrBuffer = _renderer.BufferManager.ReserveOrCreate(_renderer, cbs, hdrData.Length * sizeof(float)); + hdrBuffer.Holder.SetDataUnchecked(hdrBuffer.Offset, hdrData); + + if (hdr) + { + _pipeline.SetUniformBuffers([new BufferAssignment(2, buffer.Range), new BufferAssignment(3, hdrBuffer.Range)]); + } + else + { + _pipeline.SetUniformBuffers([new BufferAssignment(2, buffer.Range)]); + } + + _pipeline.SetImage(0, destinationTexture); + + int dispatchX = ((int)outW + (BlockWidth - 1)) / BlockWidth; + int dispatchY = ((int)outH + (BlockHeight - 1)) / BlockHeight; + + _pipeline.DispatchCompute(dispatchX, dispatchY, 1); + _pipeline.ComputeBarrier(); + + _pipeline.Finish(); + } + + // Port of NVScalerUpdateConfig (NIS_Config.h, MIT) for the SDR path. Fills the + // NISConfig constant buffer in the exact field order expected by NisScaling.glsl. + private static void BuildConfig(Span cb, float sharpness, uint inW, uint inH, uint outW, uint outH) + { + sharpness = Math.Clamp(sharpness, 0f, 1f); + float sharpenSlider = sharpness - 0.5f; // map 0..1 to -0.5..+0.5 + + float maxScale = sharpenSlider >= 0.0f ? 1.25f : 1.75f; + float minScale = sharpenSlider >= 0.0f ? 1.25f : 1.0f; + float limitScale = sharpenSlider >= 0.0f ? 1.25f : 1.0f; + + float kDetectRatio = 2f * 1127f / 1024f; + + // SDR params (HDR variant handled separately later). + float kDetectThres = 64f / 1024f; + float kMinContrastRatio = 2.0f; + float kMaxContrastRatio = 10.0f; + + float kSharpStartY = 0.45f; + float kSharpEndY = 0.9f; + float kSharpStrengthMin = MathF.Max(0.0f, 0.4f + sharpenSlider * minScale * 1.2f); + float kSharpStrengthMax = 1.6f + sharpenSlider * maxScale * 1.8f; + float kSharpLimitMin = MathF.Max(0.1f, 0.14f + sharpenSlider * limitScale * 0.32f); + float kSharpLimitMax = 0.5f + sharpenSlider * limitScale * 0.6f; + + float kRatioNorm = 1.0f / (kMaxContrastRatio - kMinContrastRatio); + float kSharpScaleY = 1.0f / (kSharpEndY - kSharpStartY); + float kSharpStrengthScale = kSharpStrengthMax - kSharpStrengthMin; + float kSharpLimitScale = kSharpLimitMax - kSharpLimitMin; + + float kSrcNormX = 1f / inW; + float kSrcNormY = 1f / inH; + float kDstNormX = 1f / outW; + float kDstNormY = 1f / outH; + float kScaleX = inW / (float)outW; + float kScaleY = inH / (float)outH; + + int i = 0; + cb[i++] = kDetectRatio; + cb[i++] = kDetectThres; + cb[i++] = kMinContrastRatio; + cb[i++] = kRatioNorm; + cb[i++] = 1.0f; // kContrastBoost + cb[i++] = 1.0f / 255.0f; // kEps + cb[i++] = kSharpStartY; + cb[i++] = kSharpScaleY; + cb[i++] = kSharpStrengthMin; + cb[i++] = kSharpStrengthScale; + cb[i++] = kSharpLimitMin; + cb[i++] = kSharpLimitScale; + cb[i++] = kScaleX; + cb[i++] = kScaleY; + cb[i++] = kDstNormX; + cb[i++] = kDstNormY; + cb[i++] = kSrcNormX; + cb[i++] = kSrcNormY; + cb[i++] = AsFloat(0); // kInputViewportOriginX + cb[i++] = AsFloat(0); // kInputViewportOriginY + cb[i++] = AsFloat(inW); // kInputViewportWidth + cb[i++] = AsFloat(inH); // kInputViewportHeight + cb[i++] = AsFloat(0); // kOutputViewportOriginX + cb[i++] = AsFloat(0); // kOutputViewportOriginY + cb[i++] = AsFloat(outW); // kOutputViewportWidth + cb[i++] = AsFloat(outH); // kOutputViewportHeight + cb[i++] = 0f; // reserved0 + cb[i++] = 0f; // reserved1 + } + + // Reinterpret a uint's bits as a float so the UBO's uint members get the right bytes. + private static float AsFloat(uint value) => BitConverter.Int32BitsToSingle((int)value); + } +} diff --git a/src/Ryujinx.Graphics.Vulkan/Effects/Shaders/FsrSharpening.glsl b/src/Ryujinx.Graphics.Vulkan/Effects/Shaders/FsrSharpening.glsl index f69a20b88..249db6c50 100644 --- a/src/Ryujinx.Graphics.Vulkan/Effects/Shaders/FsrSharpening.glsl +++ b/src/Ryujinx.Graphics.Vulkan/Effects/Shaders/FsrSharpening.glsl @@ -24,6 +24,7 @@ layout( binding = 4 ) uniform sharpening float curve_data; float gamma_data; float blend_data; + float whiten_data; #endif }; @@ -3916,8 +3917,14 @@ vec3 sdrToHdr(vec3 srgb) float luma = dot(lin, vec3(0.2126, 0.7152, 0.0722)); float maxc = max(max(lin.r, lin.g), lin.b); float hl = mix(luma, maxc, blend_data); - float boost = mix(1.0, peak / paperWhite, pow(clamp(hl, 0.0, 1.0), curve)); - return lin * paperWhite * boost; + float hlAmount = pow(clamp(hl, 0.0, 1.0), curve); + float boost = mix(1.0, peak / paperWhite, hlAmount); + vec3 outc = lin * paperWhite * boost; + // Desaturate the most extreme highlights toward white (same change as the plain blit path): + // a blinding light reads as white, not a saturated colour. Luminance-preserving. + float t = whiten_data * hlAmount; + float lumaOut = dot(outc, vec3(0.2126, 0.7152, 0.0722)); + return mix(outc, vec3(lumaOut), t); } #endif diff --git a/src/Ryujinx.Graphics.Vulkan/Effects/Shaders/FsrSharpeningHdr.spv b/src/Ryujinx.Graphics.Vulkan/Effects/Shaders/FsrSharpeningHdr.spv index 8ca5e52bb098b9d1a216af42c86e9095460db136..3b106f7d7b465bc15a8d859cfa132cc057d472e6 100644 GIT binary patch delta 3236 zcmZ9N%}-oa7{<>q!$&Psir4~03v@)Rw1!$fMQW+Q6k7{a`lYtRbQodm41+L3i?-gj zSV|Y#G}p$pJCi0RhQGkXr7PX&PUA)s6Bp{jM2ydG=AL+&NiOGkp7-;-=iJMmUj#ni z2!w-WkxF9%X1fWSdk@0V^7BDsESmEL@u?QG$v%Ke156vu-(=|lrd4w z8@^ARRB*1-cHt(Evk&+lkKguqV}S?Fs<Y z+sv?=*&T64NUPtlR=nVhgkQb*S>bA7B(k0p2neI&4`*a(*x&FlG1krqvC}1bTs&0d zPC>XLp~~lQOZ=$#0X0h@W{&nq8qZ#}!>qp%rs#*lj1Y~YaK?;=FJra7cRd|#z*L3| z!n*sgzQf_x>)TRD!OE?-{De4{-tw?GSKe~$Tx~z+YWq1?+t0b$@F88DU&_?&q&w75 z=bV;2spp6~rv-SwCc$oA7ar3}tJns$TMl?i7>yVX8BSV`<4rvr(RcF0pCfq7)A45; z@P7FpGR83%MY`mK!6u0bbB_Kb)8gn~E0Cm@h0!=)8gB^d-{2ec{l2$*>~oQc%zS_H<$3XjgKZ1Wlp8!YJF4q z*zFV!XzkEbBYa=$T|Kmhy;j;Ex+2vo1b^J~=ZjA)cG~8@J~kWogz)c*vwWBMZ{q8w z+$kj5EdeJC)~Q~YM7FsGVVZkh4^0t+4+Ysfc4DVw+avs=uV9ln1wYogV1H5^pQfPm zf5E3DkjOgi6~+;rA`BFb8=#MnX%+riYl~Lf^gdy9%Ok>U;*MgN*we!Jpp&zimkiN% z{){k%KJfzDh0*>|Ap+Wj(X9Wo!uZ>Q`-SnDP&sz`oG{vcxAkn@5O=^@4hVPn9_CEM zL1}i*WG4f@|LWs#(A}KsbvniDO=CS=!n;a%w~zDn@A2^w@jf5t<_>wBJ@K*X6`$71 z{uS#K1dM#o`(@-ij*&$0e*M$dyv&Qjr?t``P0&%(Ih=Nx*skvArUEQS}Q3 z97u#n@UYgXo>7l8y5qvrTCdCAt948q%@%uAm>iSdbvq%9#?NXeJ&j9&Hq1qS%@Y~( dga+5I^^`Oc5P+SxG@-&YMPcllx?6pE{sT2AIg9`R delta 2692 zcmZ9N%WqUw9LIn2=)=&FX~90I_BFMrgN4>Y8EX-+Fw_<)m0D0L9jew!%k(j=l9*iM zMpq={#)XM-Z9+mq`Y*6>;Yv3yHEuL9aRCc97{8yHbIhH&$<6tGzrW`>=XYm**o{59 z7fU7LnFh0%?XgDt^iC>YcQ#=*edo%-`#-j)!pEuGLDHU z11{b?A}U+bdd1fZn}s#U`)YWmhWFR-ff{~zrb?)@5%HC^7N!p1j zp0GE>mp{9u1{>{c_|4w#m3);)y-nk>dTqYAc_Y7B++LrYsBGt*D!-(?U&BAF;j2}= z-ZsML(rvMg@MyX_wIv%*oz+_@EJ!bfx6<9Qo$%ZA-SBdAf56^)D^8ZzVZ!^w350Vu zdxSe3r)&5D$JZV2DD#AEh{uJ`sV0~wQ(4OS=!kmwDF10gpde3PWlQ2d0hlLKS;`4^ zq>%s<4esoYxBQhzHrns;aIp2~fgWkyGSqxao*I@*{uBn@FFTxib*%(TAD zatc;v+Vd%KX4>Wof4kJ7rS0!xbxFsG{y92g;RvfnbgaL{X(4lxsW%@dQb^qR|q);49Pep zJfaodCpj<7twxi)C64|;aU|te1dZGFaTkS&;}U)GS#k6SfeULgZ~p2TtM`guH&fJH?ky*7H#AMaT+0xVsQL?BOf_C zHii-I8(B{1{^}znBm9i`K!neU zA9kFs_<)`lU)0Jr6&sWVjQqwNl(9!IMiRku`ft%z?FHdktu#neG(vIw{FYx7#?Gj* znU{pwAfrZ3I6mG#%Y2f|2OJV6zz;h-tJep-EKFhVKPZg (g_0_90_min * kDetectRatio)) && (g_0_90_max > kDetectThres) && (g_0_90_max > g_45_135_min); + NVB c_45_135 = (g_45_135_max > (g_45_135_min * kDetectRatio)) && (g_45_135_max > kDetectThres) && (g_45_135_max > g_0_90_min); + NVB c_g_0_90 = g_0_90_max == g_0; + NVB c_g_45_135 = g_45_135_max == g_45; + + NVF f_e_0_90 = (c_0_90 && c_45_135) ? e_0_90 : 1.0f; + NVF f_e_45_135 = (c_0_90 && c_45_135) ? e_45_135 : 1.0f; + + NVF weight_0 = (c_0_90 && c_g_0_90) ? f_e_0_90 : 0.0f; + NVF weight_90 = (c_0_90 && !c_g_0_90) ? f_e_0_90 : 0.0f; + NVF weight_45 = (c_45_135 && c_g_45_135) ? f_e_45_135 : 0.0f; + NVF weight_135 = (c_45_135 && !c_g_45_135) ? f_e_45_135 : 0.0f; + + return NVF4(weight_0, weight_90, weight_45, weight_135); +} + +#if NIS_SCALER + +#ifndef NIS_BLOCK_WIDTH +#define NIS_BLOCK_WIDTH 32 +#endif +#ifndef NIS_BLOCK_HEIGHT +#define NIS_BLOCK_HEIGHT 24 +#endif +#ifndef NIS_THREAD_GROUP_SIZE +#define NIS_THREAD_GROUP_SIZE 256 +#endif +#define kPhaseCount 64 +#define kFilterSize 6 +#define kSupportSize 6 +#define kPadSize kSupportSize +// 'Tile' is the region of source luminance values that we load into shPixelsY. +// It is the area of source pixels covered by the destination 'Block' plus a +// 3 pixel border of support pixels. +#define kTilePitch (NIS_BLOCK_WIDTH + kPadSize) +#define kTileSize (kTilePitch * (NIS_BLOCK_HEIGHT + kPadSize)) +// 'EdgeMap' is the region of source pixels for which edge map vectors are derived. +// It is the area of source pixels covered by the destination 'Block' plus a +// 1 pixel border. +#define kEdgeMapPitch (NIS_BLOCK_WIDTH + 2) +#define kEdgeMapSize (kEdgeMapPitch * (NIS_BLOCK_HEIGHT + 2)) + +NVSHARED NVF shPixelsY[kTileSize]; +NVSHARED NVH shCoefScaler[kPhaseCount][kFilterSize]; +NVSHARED NVH shCoefUSM[kPhaseCount][kFilterSize]; +NVSHARED NVH4 shEdgeMap[kEdgeMapSize]; + +void LoadFilterBanksSh(NVI i0) { + // Load up filter banks to shared memory + // The work is spread over (kPhaseCount * 2) threads + NVI i = i0; +#if( kPhaseCount * 2 > NIS_THREAD_GROUP_SIZE ) + for (; i < kPhaseCount * 2; i += NIS_THREAD_GROUP_SIZE) +#else + if (i < kPhaseCount * 2) +#endif + { + NVI phase = i >> 1; + NVI vIdx = i & 1; + + NVH4 v = NVH4(NVTEX_LOAD(coef_scaler, NVI2(vIdx, phase))); + NVI filterOffset = vIdx * 4; + shCoefScaler[phase][filterOffset + 0] = v.x; + shCoefScaler[phase][filterOffset + 1] = v.y; + if (vIdx == 0) + { + shCoefScaler[phase][2] = v.z; + shCoefScaler[phase][3] = v.w; + } + + v = NVH4(NVTEX_LOAD(coef_usm, NVI2(vIdx, phase))); + shCoefUSM[phase][filterOffset + 0] = v.x; + shCoefUSM[phase][filterOffset + 1] = v.y; + if (vIdx == 0) + { + shCoefUSM[phase][2] = v.z; + shCoefUSM[phase][3] = v.w; + } + } +} + + +NVF CalcLTI(NVF p0, NVF p1, NVF p2, NVF p3, NVF p4, NVF p5, NVI phase_index) +{ + const NVB selector = (phase_index <= kPhaseCount / 2); + NVF sel = selector ? p0 : p3; + const NVF a_min = min(min(p1, p2), sel); + const NVF a_max = max(max(p1, p2), sel); + sel = selector ? p2 : p5; + const NVF b_min = min(min(p3, p4), sel); + const NVF b_max = max(max(p3, p4), sel); + + const NVF a_cont = a_max - a_min; + const NVF b_cont = b_max - b_min; + + const NVF cont_ratio = max(a_cont, b_cont) / (min(a_cont, b_cont) + kEps); + return (1.0f - saturate((cont_ratio - kMinContrastRatio) * kRatioNorm)) * kContrastBoost; +} + +NVF4 GetInterpEdgeMap(const NVF4 edge[2][2], NVF phase_frac_x, NVF phase_frac_y) +{ + NVF4 h0 = lerp(edge[0][0], edge[0][1], phase_frac_x); + NVF4 h1 = lerp(edge[1][0], edge[1][1], phase_frac_x); + return lerp(h0, h1, phase_frac_y); +} + +NVF EvalPoly6(const NVF pxl[6], NVI phase_int) +{ + NVF y = 0.f; + { + NIS_UNROLL + for (NVI i = 0; i < 6; ++i) + { + y += shCoefScaler[phase_int][i] * pxl[i]; + } + } + NVF y_usm = 0.f; + { + NIS_UNROLL + for (NVI i = 0; i < 6; ++i) + { + y_usm += shCoefUSM[phase_int][i] * pxl[i]; + } + } + + // let's compute a piece-wise ramp based on luma + const NVF y_scale = 1.0f - saturate((y * (1.0f / NIS_SCALE_FLOAT) - kSharpStartY) * kSharpScaleY); + + // scale the ramp to sharpen as a function of luma + const NVF y_sharpness = y_scale * kSharpStrengthScale + kSharpStrengthMin; + + y_usm *= y_sharpness; + + // scale the ramp to limit USM as a function of luma + const NVF y_sharpness_limit = (y_scale * kSharpLimitScale + kSharpLimitMin) * y; + + y_usm = min(y_sharpness_limit, max(-y_sharpness_limit, y_usm)); + // reduce ringing + y_usm *= CalcLTI(pxl[0], pxl[1], pxl[2], pxl[3], pxl[4], pxl[5], phase_int); + + return y + y_usm; +} + +NVF FilterNormal(const NVF p[6][6], NVI phase_x_frac_int, NVI phase_y_frac_int) +{ + NVF h_acc = 0.0f; + NIS_UNROLL + for (NVI j = 0; j < 6; ++j) + { + NVF v_acc = 0.0f; + NIS_UNROLL + for (NVI i = 0; i < 6; ++i) + { + v_acc += p[i][j] * shCoefScaler[phase_y_frac_int][i]; + } + h_acc += v_acc * shCoefScaler[phase_x_frac_int][j]; + } + + // let's return the sum unpacked -> we can accumulate it later + return h_acc; +} + +NVF AddDirFilters(NVF p[6][6], NVF phase_x_frac, NVF phase_y_frac, NVI phase_x_frac_int, NVI phase_y_frac_int, NVF4 w) +{ + NVF f = 0; + if (w.x > 0.0f) + { + // 0 deg filter + NVF interp0Deg[6]; + { + NIS_UNROLL + for (NVI i = 0; i < 6; ++i) + { + interp0Deg[i] = lerp(p[i][2], p[i][3], phase_x_frac); + } + } + f += EvalPoly6(interp0Deg, phase_y_frac_int) * w.x; + } + if (w.y > 0.0f) + { + // 90 deg filter + NVF interp90Deg[6]; + { + NIS_UNROLL + for (NVI i = 0; i < 6; ++i) + { + interp90Deg[i] = lerp(p[2][i], p[3][i], phase_y_frac); + } + } + + f += EvalPoly6(interp90Deg, phase_x_frac_int) * w.y; + } + if (w.z > 0.0f) + { + //45 deg filter + NVF pphase_b45 = 0.5f + 0.5f * (phase_x_frac - phase_y_frac); + + NVF temp_interp45Deg[7]; + temp_interp45Deg[1] = lerp(p[2][1], p[1][2], pphase_b45); + temp_interp45Deg[3] = lerp(p[3][2], p[2][3], pphase_b45); + temp_interp45Deg[5] = lerp(p[4][3], p[3][4], pphase_b45); + { + pphase_b45 = pphase_b45 - 0.5f; + NVF a = (pphase_b45 >= 0.f) ? p[0][2] : p[2][0]; + NVF b = (pphase_b45 >= 0.f) ? p[1][3] : p[3][1]; + NVF c = (pphase_b45 >= 0.f) ? p[2][4] : p[4][2]; + NVF d = (pphase_b45 >= 0.f) ? p[3][5] : p[5][3]; + temp_interp45Deg[0] = lerp(p[1][1], a, abs(pphase_b45)); + temp_interp45Deg[2] = lerp(p[2][2], b, abs(pphase_b45)); + temp_interp45Deg[4] = lerp(p[3][3], c, abs(pphase_b45)); + temp_interp45Deg[6] = lerp(p[4][4], d, abs(pphase_b45)); + } + + NVF interp45Deg[6]; + NVF pphase_p45 = phase_x_frac + phase_y_frac; + if (pphase_p45 >= 1) + { + NIS_UNROLL + for (NVI i = 0; i < 6; i++) + { + interp45Deg[i] = temp_interp45Deg[i + 1]; + } + pphase_p45 = pphase_p45 - 1; + } + else + { + NIS_UNROLL + for (NVI i = 0; i < 6; i++) + { + interp45Deg[i] = temp_interp45Deg[i]; + } + } + + f += EvalPoly6(interp45Deg, NVI(pphase_p45 * 64)) * w.z; + } + if (w.w > 0.0f) + { + //135 deg filter + NVF pphase_b135 = 0.5f * (phase_x_frac + phase_y_frac); + + NVF temp_interp135Deg[7]; + temp_interp135Deg[1] = lerp(p[3][1], p[4][2], pphase_b135); + temp_interp135Deg[3] = lerp(p[2][2], p[3][3], pphase_b135); + temp_interp135Deg[5] = lerp(p[1][3], p[2][4], pphase_b135); + { + pphase_b135 = pphase_b135 - 0.5f; + NVF a = (pphase_b135 >= 0.f) ? p[5][2] : p[3][0]; + NVF b = (pphase_b135 >= 0.f) ? p[4][3] : p[2][1]; + NVF c = (pphase_b135 >= 0.f) ? p[3][4] : p[1][2]; + NVF d = (pphase_b135 >= 0.f) ? p[2][5] : p[0][3]; + temp_interp135Deg[0] = lerp(p[4][1], a, abs(pphase_b135)); + temp_interp135Deg[2] = lerp(p[3][2], b, abs(pphase_b135)); + temp_interp135Deg[4] = lerp(p[2][3], c, abs(pphase_b135)); + temp_interp135Deg[6] = lerp(p[1][4], d, abs(pphase_b135)); + } + + NVF interp135Deg[6]; + NVF pphase_p135 = 1 + (phase_x_frac - phase_y_frac); + if (pphase_p135 >= 1) + { + NIS_UNROLL + for (NVI i = 0; i < 6; ++i) + { + interp135Deg[i] = temp_interp135Deg[i + 1]; + } + pphase_p135 = pphase_p135 - 1; + } + else + { + NIS_UNROLL + for (NVI i = 0; i < 6; ++i) + { + interp135Deg[i] = temp_interp135Deg[i]; + } + } + + f += EvalPoly6(interp135Deg, NVI(pphase_p135 * 64)) * w.w; + } + return f; +} + + +//----------------------------------------------------------------------------------------------- +// NVScaler +//----------------------------------------------------------------------------------------------- +void NVScaler(NVU2 blockIdx, NVU threadIdx) +{ + // Figure out the range of pixels from input image that would be needed to be loaded for this thread-block + NVI dstBlockX = NVI(NIS_BLOCK_WIDTH * blockIdx.x); + NVI dstBlockY = NVI(NIS_BLOCK_HEIGHT * blockIdx.y); + + const NVI srcBlockStartX = NVI(floor((dstBlockX + 0.5f) * kScaleX - 0.5f)); + const NVI srcBlockStartY = NVI(floor((dstBlockY + 0.5f) * kScaleY - 0.5f)); + const NVI srcBlockEndX = NVI(ceil((dstBlockX + NIS_BLOCK_WIDTH + 0.5f) * kScaleX - 0.5f)); + const NVI srcBlockEndY = NVI(ceil((dstBlockY + NIS_BLOCK_HEIGHT + 0.5f) * kScaleY - 0.5f)); + + NVI numTilePixelsX = srcBlockEndX - srcBlockStartX + kSupportSize - 1; + NVI numTilePixelsY = srcBlockEndY - srcBlockStartY + kSupportSize - 1; + + // round-up load region to even size since we're loading in 2x2 batches + numTilePixelsX += numTilePixelsX & 0x1; + numTilePixelsY += numTilePixelsY & 0x1; + const NVI numTilePixels = numTilePixelsX * numTilePixelsY; + + // calculate the equivalent values for the edge map + const NVI numEdgeMapPixelsX = numTilePixelsX - kSupportSize + 2; + const NVI numEdgeMapPixelsY = numTilePixelsY - kSupportSize + 2; + const NVI numEdgeMapPixels = numEdgeMapPixelsX * numEdgeMapPixelsY; + + // fill in input luma tile (shPixelsY) in batches of 2x2 pixels + // we use texture gather to get extra support necessary + // to compute 2x2 edge map outputs too + { + for (NVU i = threadIdx * 2; i < NVU(numTilePixels) >> 1; i += NIS_THREAD_GROUP_SIZE * 2) + { + NVU py = (i / numTilePixelsX) * 2; + NVU px = i % numTilePixelsX; + + // 0.5 to be in the center of texel + // - (kSupportSize - 1) / 2 to shift by the kernel support size + NVF kShift = 0.5f - (kSupportSize - 1) / 2; +#if NIS_VIEWPORT_SUPPORT + const NVF tx = (srcBlockStartX + px + kInputViewportOriginX + kShift) * kSrcNormX; + const NVF ty = (srcBlockStartY + py + kInputViewportOriginY + kShift) * kSrcNormY; +#else + const NVF tx = (srcBlockStartX + px + kShift) * kSrcNormX; + const NVF ty = (srcBlockStartY + py + kShift) * kSrcNormY; +#endif + NVF p[2][2]; +#if NIS_TEXTURE_GATHER + { + const NVF4 sr = NVTEX_SAMPLE_RED(in_texture, samplerLinearClamp, NVF2(tx, ty)); + const NVF4 sg = NVTEX_SAMPLE_GREEN(in_texture, samplerLinearClamp, NVF2(tx, ty)); + const NVF4 sb = NVTEX_SAMPLE_BLUE(in_texture, samplerLinearClamp, NVF2(tx, ty)); + + p[0][0] = getY(NVF3(sr.w, sg.w, sb.w)); + p[0][1] = getY(NVF3(sr.z, sg.z, sb.z)); + p[1][0] = getY(NVF3(sr.x, sg.x, sb.x)); + p[1][1] = getY(NVF3(sr.y, sg.y, sb.y)); + } +#else + NIS_UNROLL_INNER + for (NVI j = 0; j < 2; j++) + { + NIS_UNROLL_INNER + for (NVI k = 0; k < 2; k++) + { +#if NIS_NV12_SUPPORT + p[j][k] = NVTEX_SAMPLE(in_texture_y, samplerLinearClamp, NVF2(tx + k * kSrcNormX, ty + j * kSrcNormY)); +#else + const NVF4 px = NVTEX_SAMPLE(in_texture, samplerLinearClamp, NVF2(tx + k * kSrcNormX, ty + j * kSrcNormY)); + p[j][k] = getY(px.xyz); +#endif + } + } +#endif + const NVU idx = py * kTilePitch + px; + shPixelsY[idx] = NVH(p[0][0]); + shPixelsY[idx + 1] = NVH(p[0][1]); + shPixelsY[idx + kTilePitch] = NVH(p[1][0]); + shPixelsY[idx + kTilePitch + 1] = NVH(p[1][1]); + } + } + GroupMemoryBarrierWithGroupSync(); + { + // fill in the edge map of 2x2 pixels + for (NVU i = threadIdx * 2; i < NVU(numEdgeMapPixels) >> 1; i += NIS_THREAD_GROUP_SIZE * 2) + { + NVU py = (i / numEdgeMapPixelsX) * 2; + NVU px = i % numEdgeMapPixelsX; + + const NVU edgeMapIdx = py * kEdgeMapPitch + px; + + NVU tileCornerIdx = (py + 1) * kTilePitch + px + 1; + NVF p[4][4]; + NIS_UNROLL_INNER + for (NVI j = 0; j < 4; j++) + { + NIS_UNROLL_INNER + for (NVI k = 0; k < 4; k++) + { + p[j][k] = shPixelsY[tileCornerIdx + j * kTilePitch + k]; + } + } + + shEdgeMap[edgeMapIdx] = NVH4(GetEdgeMap(p, 0, 0)); + shEdgeMap[edgeMapIdx + 1] = NVH4(GetEdgeMap(p, 0, 1)); + shEdgeMap[edgeMapIdx + kEdgeMapPitch] = NVH4(GetEdgeMap(p, 1, 0)); + shEdgeMap[edgeMapIdx + kEdgeMapPitch + 1] = NVH4(GetEdgeMap(p, 1, 1)); + } + } + LoadFilterBanksSh(NVI(threadIdx)); + GroupMemoryBarrierWithGroupSync(); + + // output coord within a tile + const NVI2 pos = NVI2(NVU(threadIdx) % NVU(NIS_BLOCK_WIDTH), NVU(threadIdx) / NVU(NIS_BLOCK_WIDTH)); + // x coord inside the output image + const NVI dstX = dstBlockX + pos.x; + // x coord inside the input image + const NVF srcX = (0.5f + dstX) * kScaleX - 0.5f; + // nearest integer part + const NVI px = NVI(floor(srcX) - srcBlockStartX); + // fractional part + const NVF fx = srcX - floor(srcX); + // discretized phase + const NVI fx_int = NVI(fx * kPhaseCount); +#if NIS_VIEWPORT_SUPPORT + if (NVU(srcX) > kInputViewportWidth || NVU(dstX) > kOutputViewportWidth) + { + return; + } +#endif + for (NVI k = 0; k < NIS_BLOCK_WIDTH * NIS_BLOCK_HEIGHT / NIS_THREAD_GROUP_SIZE; ++k) + { + // y coord inside the output image + const NVI dstY = dstBlockY + pos.y + k * (NIS_THREAD_GROUP_SIZE / NIS_BLOCK_WIDTH); + // y coord inside the input image + const NVF srcY = (0.5f + dstY) * kScaleY - 0.5f; +#if NIS_VIEWPORT_SUPPORT + if (!(NVU(srcY) > kInputViewportHeight || NVU(dstY) > kOutputViewportHeight)) +#endif + { + // nearest integer part + const NVI py = NVI(floor(srcY) - srcBlockStartY); + // fractional part + const NVF fy = srcY - floor(srcY); + // discretized phase + const NVI fy_int = NVI(fy * kPhaseCount); + + // generate weights for directional filters + const NVI startEdgeMapIdx = py * kEdgeMapPitch + px; + NVF4 edge[2][2]; + NIS_UNROLL + for (NVI i = 0; i < 2; i++) + { + NIS_UNROLL + for (NVI j = 0; j < 2; j++) + { + // need to shift edge map sampling since it's a 2x2 centered inside 6x6 grid + edge[i][j] = shEdgeMap[startEdgeMapIdx + (i * kEdgeMapPitch) + j]; + } + } + const NVF4 w = GetInterpEdgeMap(edge, fx, fy) * NIS_SCALE_INT; + + // load 6x6 support to regs + const NVI startTileIdx = py * kTilePitch + px; + NVF p[6][6]; + { + NIS_UNROLL + for (NVI i = 0; i < 6; ++i) + { + NIS_UNROLL + for (NVI j = 0; j < 6; ++j) + { + p[i][j] = shPixelsY[startTileIdx + i * kTilePitch + j]; + } + } + } + + // weigth for luma + const NVF baseWeight = NIS_SCALE_FLOAT - w.x - w.y - w.z - w.w; + + // final luma is a weighted product of directional & normal filters + NVF opY = 0; + + // get traditional scaler filter output + opY += FilterNormal(p, fx_int, fy_int) * baseWeight; + + // get directional filter bank output + opY += AddDirFilters(p, fx, fy, fx_int, fy_int, w); + +#if NIS_VIEWPORT_SUPPORT + NVF2 coord = NVF2((srcX + kInputViewportOriginX + 0.5f) * kSrcNormX, (srcY + kInputViewportOriginY + 0.5f) * kSrcNormY); + NVF2 dstCoord = NVF2(dstX + kOutputViewportOriginX, dstY + kOutputViewportOriginY); +#else + NVF2 coord = NVF2((srcX + 0.5f) * kSrcNormX, (srcY + 0.5f) * kSrcNormY); + NVF2 dstCoord = NVF2(dstX, dstY); +#endif + // do bilinear tap for chroma upscaling +#if NIS_NV12_SUPPORT + NVF y = NVTEX_SAMPLE(in_texture_y, samplerLinearClamp, coord); + NVF2 uv = NVTEX_SAMPLE(in_texture_uv, samplerLinearClamp, coord); + NVF4 op = NVF4(YUVtoRGB(NVF3(y, uv)), 1.0f); +#else + NVF4 op = NVTEX_SAMPLE(in_texture, samplerLinearClamp, coord); + NVF y = getY(NVF3(op.x, op.y, op.z)); +#endif + +#if NIS_HDR_MODE == NIS_HDR_MODE_LINEAR + const NVF kEps = 1e-4f; + const NVF kNorm = 1.0f / (NIS_SCALE_FLOAT * kHDRCompressionFactor); + const NVF opYN = max(opY, 0.0f) * kNorm; + const NVF corr = (opYN * opYN + kEps) / (max(getYLinear(NVF3(op.x, op.y, op.z)), 0.0f) + kEps); + op.x *= corr; + op.y *= corr; + op.z *= corr; +#else + const NVF corr = opY * (1.0f / NIS_SCALE_FLOAT) - y; + op.x += corr; + op.y += corr; + op.z += corr; +#endif + NVTEX_STORE(out_texture, dstCoord, NVCLAMP(op)); + } + } +} +#else + +#ifndef NIS_BLOCK_WIDTH +#define NIS_BLOCK_WIDTH 32 +#endif +#ifndef NIS_BLOCK_HEIGHT +#define NIS_BLOCK_HEIGHT 32 +#endif +#ifndef NIS_THREAD_GROUP_SIZE +#define NIS_THREAD_GROUP_SIZE 256 +#endif + +#define kSupportSize 5 +#define kNumPixelsX (NIS_BLOCK_WIDTH + kSupportSize + 1) +#define kNumPixelsY (NIS_BLOCK_HEIGHT + kSupportSize + 1) + +NVSHARED NVF shPixelsY[kNumPixelsY][kNumPixelsX]; + +NVF CalcLTIFast(const NVF y[5]) +{ + const NVF a_min = min(min(y[0], y[1]), y[2]); + const NVF a_max = max(max(y[0], y[1]), y[2]); + + const NVF b_min = min(min(y[2], y[3]), y[4]); + const NVF b_max = max(max(y[2], y[3]), y[4]); + + const NVF a_cont = a_max - a_min; + const NVF b_cont = b_max - b_min; + + const NVF cont_ratio = max(a_cont, b_cont) / (min(a_cont, b_cont) + kEps); + return (1.0f - saturate((cont_ratio - kMinContrastRatio) * kRatioNorm)) * kContrastBoost; +} + +NVF EvalUSM(const NVF pxl[5], const NVF sharpnessStrength, const NVF sharpnessLimit) +{ + // USM profile + NVF y_usm = -0.6001f * pxl[1] + 1.2002f * pxl[2] - 0.6001f * pxl[3]; + // boost USM profile + y_usm *= sharpnessStrength; + // clamp to the limit + y_usm = min(sharpnessLimit, max(-sharpnessLimit, y_usm)); + // reduce ringing + y_usm *= CalcLTIFast(pxl); + + return y_usm; +} + +NVF4 GetDirUSM(const NVF p[5][5]) +{ + // sharpness boost & limit are the same for all directions + const NVF scaleY = 1.0f - saturate((p[2][2] - kSharpStartY) * kSharpScaleY); + // scale the ramp to sharpen as a function of luma + const NVF sharpnessStrength = scaleY * kSharpStrengthScale + kSharpStrengthMin; + // scale the ramp to limit USM as a function of luma + const NVF sharpnessLimit = (scaleY * kSharpLimitScale + kSharpLimitMin) * p[2][2]; + + NVF4 rval; + // 0 deg filter + NVF interp0Deg[5]; + { + for (NVI i = 0; i < 5; ++i) + { + interp0Deg[i] = p[i][2]; + } + } + + rval.x = EvalUSM(interp0Deg, sharpnessStrength, sharpnessLimit); + + // 90 deg filter + NVF interp90Deg[5]; + { + for (NVI i = 0; i < 5; ++i) + { + interp90Deg[i] = p[2][i]; + } + } + + rval.y = EvalUSM(interp90Deg, sharpnessStrength, sharpnessLimit); + + //45 deg filter + NVF interp45Deg[5]; + interp45Deg[0] = p[1][1]; + interp45Deg[1] = lerp(p[2][1], p[1][2], 0.5f); + interp45Deg[2] = p[2][2]; + interp45Deg[3] = lerp(p[3][2], p[2][3], 0.5f); + interp45Deg[4] = p[3][3]; + + rval.z = EvalUSM(interp45Deg, sharpnessStrength, sharpnessLimit); + + //135 deg filter + NVF interp135Deg[5]; + interp135Deg[0] = p[3][1]; + interp135Deg[1] = lerp(p[3][2], p[2][1], 0.5f); + interp135Deg[2] = p[2][2]; + interp135Deg[3] = lerp(p[2][3], p[1][2], 0.5f); + interp135Deg[4] = p[1][3]; + + rval.w = EvalUSM(interp135Deg, sharpnessStrength, sharpnessLimit); + return rval; +} + +//----------------------------------------------------------------------------------------------- +// NVSharpen +//----------------------------------------------------------------------------------------------- +void NVSharpen(NVU2 blockIdx, NVU threadIdx) +{ + const NVI dstBlockX = NVI(NIS_BLOCK_WIDTH * blockIdx.x); + const NVI dstBlockY = NVI(NIS_BLOCK_HEIGHT * blockIdx.y); + + // fill in input luma tile in batches of 2x2 pixels + // we use texture gather to get extra support necessary + // to compute 2x2 edge map outputs too + const NVF kShift = 0.5f - kSupportSize / 2; + + for (NVI i = NVI(threadIdx) * 2; i < kNumPixelsX * kNumPixelsY / 2; i += NIS_THREAD_GROUP_SIZE * 2) + { + NVU2 pos = NVU2(NVU(i) % NVU(kNumPixelsX), NVU(i) / NVU(kNumPixelsX) * 2); + NIS_UNROLL + for (NVI dy = 0; dy < 2; dy++) + { + NIS_UNROLL + for (NVI dx = 0; dx < 2; dx++) + { +#if NIS_VIEWPORT_SUPPORT + const NVF tx = (dstBlockX + pos.x + kInputViewportOriginX + dx + kShift) * kSrcNormX; + const NVF ty = (dstBlockY + pos.y + kInputViewportOriginY + dy + kShift) * kSrcNormY; +#else + const NVF tx = (dstBlockX + pos.x + dx + kShift) * kSrcNormX; + const NVF ty = (dstBlockY + pos.y + dy + kShift) * kSrcNormY; +#endif +#if NIS_NV12_SUPPORT + shPixelsY[pos.y + dy][pos.x + dx] = NVTEX_SAMPLE(in_texture_y, samplerLinearClamp, NVF2(tx, ty)); +#else + const NVF4 px = NVTEX_SAMPLE(in_texture, samplerLinearClamp, NVF2(tx, ty)); + shPixelsY[pos.y + dy][pos.x + dx] = getY(px.xyz); +#endif + } + } + } + + GroupMemoryBarrierWithGroupSync(); + + for (NVI k = NVI(threadIdx); k < NIS_BLOCK_WIDTH * NIS_BLOCK_HEIGHT; k += NIS_THREAD_GROUP_SIZE) + { + const NVI2 pos = NVI2(NVU(k) % NVU(NIS_BLOCK_WIDTH), NVU(k) / NVU(NIS_BLOCK_WIDTH)); + + // load 5x5 support to regs + NVF p[5][5]; + NIS_UNROLL + for (NVI i = 0; i < 5; ++i) + { + NIS_UNROLL + for (NVI j = 0; j < 5; ++j) + { + p[i][j] = shPixelsY[pos.y + i][pos.x + j]; + } + } + + // get directional filter bank output + NVF4 dirUSM = GetDirUSM(p); + + // generate weights for directional filters + NVF4 w = GetEdgeMap(p, kSupportSize / 2 - 1, kSupportSize / 2 - 1); + + // final USM is a weighted sum filter outputs + const NVF usmY = (dirUSM.x * w.x + dirUSM.y * w.y + dirUSM.z * w.z + dirUSM.w * w.w); + + // do bilinear tap and correct rgb texel so it produces new sharpened luma + const NVI dstX = dstBlockX + pos.x; + const NVI dstY = dstBlockY + pos.y; + +#if NIS_VIEWPORT_SUPPORT + NVF2 coord = NVF2((dstX + kInputViewportOriginX + 0.5f) * kSrcNormX, (dstY + kInputViewportOriginY + 0.5f) * kSrcNormY); + NVF2 dstCoord = NVF2(dstX + kOutputViewportOriginX, dstY + kOutputViewportOriginY); + if (!(NVU(dstX) > kOutputViewportWidth || NVU(dstY) > kOutputViewportHeight)) +#else + NVF2 coord = NVF2((dstX + 0.5f) * kSrcNormX, (dstY + 0.5f) * kSrcNormY); + NVF2 dstCoord = NVF2(dstX, dstY); +#endif + { +#if NIS_NV12_SUPPORT + NVF y = NVTEX_SAMPLE(in_texture_y, samplerLinearClamp, coord); + NVF2 uv = NVTEX_SAMPLE(in_texture_uv, samplerLinearClamp, coord); + NVF4 op = NVF4(YUVtoRGB(NVF3(y, uv)), 1.0f); +#else + NVF4 op = NVTEX_SAMPLE(in_texture, samplerLinearClamp, coord); +#endif +#if NIS_HDR_MODE == NIS_HDR_MODE_LINEAR + const NVF kEps = 1e-4f * kHDRCompressionFactor * kHDRCompressionFactor; + NVF newY = p[2][2] + usmY; + newY = max(newY, 0.0f); + const NVF oldY = p[2][2]; + const NVF corr = (newY * newY + kEps) / (oldY * oldY + kEps); + op.x *= corr; + op.y *= corr; + op.z *= corr; +#else + op.x += usmY; + op.y += usmY; + op.z += usmY; +#endif + NVTEX_STORE(out_texture, dstCoord, NVCLAMP(op)); + } + } +} +#endif \ No newline at end of file diff --git a/src/Ryujinx.Graphics.Vulkan/Effects/Shaders/NisScaling.glsl b/src/Ryujinx.Graphics.Vulkan/Effects/Shaders/NisScaling.glsl new file mode 100644 index 000000000..1aae48625 --- /dev/null +++ b/src/Ryujinx.Graphics.Vulkan/Effects/Shaders/NisScaling.glsl @@ -0,0 +1,88 @@ +// The MIT License (MIT) +// +// Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. 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. +// +// NVIDIA Image Scaling (NIS) v1.0.3 - MIT License, (c) 2022 NVIDIA CORPORATION. +// Vulkan compute wrapper adapted for Ryujinx: combined samplers + inline coefficient +// tables (see NisScaler.h macro edits and nis_coef.glsl). Output path: SDR (rgba8). + +#version 450 core +#extension GL_GOOGLE_include_directive : require + +#define NIS_GLSL 1 +#define NIS_SCALER 1 +#define NIS_THREAD_GROUP_SIZE 256 +#define NIS_BLOCK_WIDTH 32 +#define NIS_BLOCK_HEIGHT 24 + +layout(local_size_x = NIS_THREAD_GROUP_SIZE, local_size_y = 1, local_size_z = 1) in; + +// NISConfig constant buffer (exact field order; filled host-side by NisScalingFilter.cs). +layout(binding = 2) uniform cb +{ + float kDetectRatio; + float kDetectThres; + float kMinContrastRatio; + float kRatioNorm; + + float kContrastBoost; + float kEps; + float kSharpStartY; + float kSharpScaleY; + + float kSharpStrengthMin; + float kSharpStrengthScale; + float kSharpLimitMin; + float kSharpLimitScale; + + float kScaleX; + float kScaleY; + + float kDstNormX; + float kDstNormY; + float kSrcNormX; + float kSrcNormY; + + uint kInputViewportOriginX; + uint kInputViewportOriginY; + uint kInputViewportWidth; + uint kInputViewportHeight; + + uint kOutputViewportOriginX; + uint kOutputViewportOriginY; + uint kOutputViewportWidth; + uint kOutputViewportHeight; + + float reserved0; + float reserved1; +}; + +layout(binding = 1, set = 2) uniform sampler2D in_texture; +layout(rgba8, binding = 0, set = 3) uniform image2D out_texture; + +// coef_scaler[64][2] and coef_usm[64][2] (vec4 = taps0-3 / taps4-7), MIT, verbatim values. +#include "nis_coef.glsl" + +#include "NisScaler.h" + +void main() +{ + NVScaler(uvec2(gl_WorkGroupID.xy), gl_LocalInvocationID.x); +} diff --git a/src/Ryujinx.Graphics.Vulkan/Effects/Shaders/NisScaling.spv b/src/Ryujinx.Graphics.Vulkan/Effects/Shaders/NisScaling.spv new file mode 100644 index 0000000000000000000000000000000000000000..b296cd6d88295c260158322b858642a22c59e4e0 GIT binary patch literal 56760 zcmb`wb-Z2G*@e495`w!0O>oyBAwUR(V8LBOJb@Sy2u^T!3KT8w#f!Up@fIz#IE7NA z_UJ^KkhEGbByuIeCJ$q$=>_qoQ~5h)}_(t(3r6?UE|{E8_n13jcHIC zXwx?Pjr(ucf5izChpgCR^_6v)rO~nJr_Zd7E{)FgooJ&6jvS-mO!EAWHXHHkvn%4P zzpW$xn}+^jNTJ=(Z#;D3 zMni@V-D2RlrH8GuTaP{Xuh+;`{7;{bjd}1H*Yxk$n3sNJ3(rTtcN0(7Scty=*nva( zjT|*`==co=j@f6zb|aQ6zsKNcW-;@D+)&83r(;dHY&6tjj zMa8SM)w7s*)i%7ic(pdXgt$i=UQ)cexN~DE@rZ#Fh7K4xX2{UVC^I%l<-9i?L$%{F z`~8OX*lpE4_^(wHKCan{^g~%7I{37WmFU~@4I4jj@PNtD-5iB}K%3sBu@+#X{RWQO zdhDnJ)^Lf}Xs$20IyTlYX7VU-=f;{<{KzpA!LC(r`nEOQa_sog14p%mw=%RXRWr{z z#x`TRG}aY2=R0{ov%=VZyRyK)WO=05##Y6tJJGG zW?ysBT$7EeylrDbn|i-WZyT>o@6y<~(*L){O)6Wx#*U3mvF%TDF5A*?x#M<&2aXy# ze(C)tt=emnJ62og#@6(MMvWc3&!$6K$ER~+8?=cd#t$7hq)l}FhoKE0urd*y8fP_e z4|a_?IyD|>@}3^nCjYd_S6NL5$C#V*m)HS4R~|5WAfu>tEi|tNuFXPCt;gyE2yOFh z)4Ttd-gfP9+$GS5HuJXST&l_2;!{g6w8526$Hsc{K@CRZ7=wN7(|71Zt{>YDoH%l< zJT}cO+x9HSgf`oB%{Z@1`)o0C%!XsfOdLOOLKQqM(dPBaw@)*p`*CVcuRZ%DXoIn1 zC$NF1W%kB&*0#~O);MkTO{=!uh=JqBZ8ve?_=&qftNOSg$L6tZ+1l!xZmMnK_@QHl zPaHw9&Db_Om)*p+?9KgA<@Wm5Hg5lsqeo6`22abi)jGXi7Phv)X709m%+$17($3Y~ zX0Ehd(Aw5$7PY<;Cb~j9H{;v%Sqt4s#eX*S?ZywT`+MJr-sV4jv(7oxHytx>(!?D{ z4&8s;*zps$8b5OQ$T4kWwfW5XfB3ZJX`Am{|KqdM$RQI)P*+=yx&Oz$$G|pseNm_^_pij^A(S5Kr9b3lzH7 ze$V~&JrSCiYXf-jfbtyc+!)-{+RpI!4R2}j^qk)ZZP+p%$Zn?{ySjcdRI2aIl>!}342aWAJ$`(5xsh5dmx zUfCaO`VZ!6)#m?H(;nWr@f=(>V0_zE&~^I=Z3IU^{DqheA^s73Ky%!Vji116Ys+1j z3)p|pwNqmbv;ziAn!qKn89S#wP29QBtsVcRoBk8Lc;Mf;(W8xf0UR@Q!UU8qjXwXa z@eIfnspG!G_Y0MY+!Do%{mtZ zXYNabhctDbE5NP%Bx>{B87=eQOI!@&oJUpmPL0Ejua0Bq#!2#4y`I?CdpcTCFLGPK z=DG%3>b)JDdLII(-sixn_gir4ZOh-W@he(ey<8o+*|ybN4bzqGpjmyD2Z;Tw)ymobLTnV4pv~_IUB5uXp4{kg6 zIyN3RrdfB##`EB|T%8*)fk%!RFmdSQiIc_;Z5!)j?M>fKjnBl@zU$cdMt>fSJ$Ao< zCyZ#DciY@}pT~n%>R%k3`j-_qYwX-u3GDkK-Z(bzG#upR?CN z&m8)KGl%}*wi-J(27xDxXdQ-D-9xljb$4nEtMTDAKDx%o*7&#@pH$=f*ZAZbKeWaV zukk50e%ycYj*S!LW4+#YYMcq?JYv_j?T2&WoG&fjv2lsMt^IHncvus6YFyvMUd`H$ z^G)!!_3YHR6HHDIV{?oL;pA-bj*TbbIhHSghc)+Dr^YKyoSf4(-Z0K5^VZ?;NV$MQ^W+cM3uWoJ9m*)I?>09idb!)tDjrXhZO=^6z8t-4@Th{pY z@WErpj?eXYFF40%!6)KMq4yul6Ue4x_M_uX;TWxLW1UkQ>)e8GOkKV| z^1b0EWE1;ZtgWjwWwlPU_!*#Md` zp>=I(Lki9JTl$SCG~a8fjiPPLZt#6pcd&aQd{MA3x%V#ZxK&(x{N3l;ed#aP-m%5q z^Byy6(e7uja_f*^)Z*4ZCt1zgooM65^MKq}-DvaT?|zbN?*?}-soA>2N!PTi&kc7U zbffuNuEpI4?&$^je~-WU)Z`v}HM#jc?&elEpU1o#twXc#M9pn>b?ft3o6r2+K^|u{ z^XO}{PUr9Z84HHvGrM-Uf$-zGbPw~b$x8ktq;9(QyZ~g?#|{+_%qMSG&i2IL%l1IB@*6Po`N%@-L~8*IcN}ftP0%uER=TZDmW$8qJu) zdppmy!Oq**h54V?d*}84M|17|-&)47j^2#xyd1w*i?83(wD)OYW3o2qqaEJjj&mh^ zeR}uORbcn2TGr%haMr^$5HHWM)$baxcGoTIbS?Ocnb+$TejWUpIk&zy{CfC{eRp0b z`EP(5qo4V;$M;6?@h|OpPkjFfS4%x*EjOW?OP|Ew4DL1V`Kp#%;Pd?SOjXOR@LMiC zthvuzquby^_IRtR?RL01jMG_``MZZ!wcQK<=%;1Y&05?CH-~ZhYESI_;2j=Xx|-($@Ue4rTsyH3!u3&4 zy*Z{2p_yBsMX2LEA|VO6b-Ze}-p*%KwLK3vhjIF9 zPwp4M6Z@T2)%GG>t-ad*jBYM{68{qTuAh&sYI_;parzUg+FpTQ*5Tl)wpZbg{d!AP z+iP%h7^koH>Ni{$M`RBebnp6_^)W@)+hPi1a}#5ZdL1B@CRnO zpsMw4_>YGkSk?Lt{F`51scL-}ZVu!0)t=n%fnV(VMpf(I;8&jiXjSX`aDCM4wSItR zZheyP@8FBRysE18AMk_by0NPDL-;$>98lHz5q!~KzN>2e7;X;Z^wpl+pMW>n`sb?F zPvM+9y|O<)gX^PSuk~{@bL*3QUw}W~=FY0tf5N|*`TnZbFX4;+vR_r}SMV*4nt8p9 z_b>fv#_6j)xxWT;4XJYd8@}zP3stp#1J_5rUhB7L=GG_qz5{>x)FV}`-@^xVdA`d3 zANZXMPpWGD0Y3c+i&VA#2seju`f5+^pTHNaw&;DC*UxaZYiae`enB&rK8gPoe9fFM zR<->GKkl11tJ=H~c6eueRa*!6AO5pSRogUha~Rk3g(r7M@F6d*R^{ph=e|>|V`sQN z>h)TuMKiZP$=3zUey?hs4*tWepH%s$hi^G(Tvh7~@UxEHsH$~FxH*i|S9@~L1U};Z zjjLQU!{1xsw5rxw;QFZ7Yn>I%-1;QnY~Yvg{h+FKcKEHI{=3RQ2mGjuMpw1Y3E%qX zovT{sf}6uQeYGd|+~DKB*{#Yo4}8VbE~#pr7p{+bz1I2A%&kxIbp`+U#b;Hm^TXf# z;nynv0`Mar+`FoEL3rO2M^&{h1UH9q`f5+^g~5G)99`vF1pf5)4^_2xgX^PSueCdx zx%ElDMZu>``&CuzV(_t3I^7@p;_xHhA70hE1bmJo53g!n5^fIT^wpl&rR392sm`UP z;c7ms)N5M?-CX)4ep&GCYyPLIZ8`X}hfP=2wmf|7lp$4ZE5MI@@Pw+i72)PEPG9ZG zy%KooSx>Ciab@`HJ-)2gaTT~e>h)SzMKiZP$+sGq>v2_U5BNMk%~;jCI{cj9_Nr=K z1Aa*Fv#VOygqy=SeYGd|TH^W6u4?TGzw75N)$`xlaDCM4we~_Yw?4_Y4tTHqX02*l z7k<=5dsMZp2jBnxi>umt!>!FYeYGdo`rrpYzo@FM54_8O`KsDBfa{}PuWds#bL*3Q zeZgaTG}aB@2!3zpIjUOw!9Ra+m#Wr{;XfXFZB^?gaB~=^ulD5H6#U4yS68)e2Iu*+ zs&#X?KI-*a`=gm#pXA#D{KEHhR<&&jfBdZ-s@k@KAJY5ws$UBUW^R3wZx8XY^HjC%3IA#LEvniEz-OA~nX0zE z;MQiGzS@&(Ao$7|9sCv zB=}?9)~{;Y4}S4FpI5c*4>!iz^wpkRlhwPvU)?7TfUEhQqF&2^=;qQV@dtsYTenv= z_QCM$*89F1`w;k|H206b+7o*y_|=a-t=8o*IPayZbvYcak9xh9DQM=_C;5&5&((X4 zs+J?+-THLso#)A;;9Y5+XZmVS?9t$#=bQRWa}4~0E6%TKITo&udcBt8(9Eq*@*NNU z$G}ypT26p3_3*S+EhoZfqj^2jS9@Yl0^dIOPu04d44>ik%c@#Vf$O7QujN!UbL*3Q zr-9eM{M~B5oDP5fxT)9mGvFf@UB0UIO!x{L&Q#TU7Tg@h>8m}t&jv3sZHMZd{T=+V zZZ}o6o&(oMyzf&>{MW*lK6JsV zw(H===%>w=b-BK!Yp;)e1Ddh=#Q#RHby}};bgr@e5$yb3hkCA?(2dn6{x^d?9v&Bu zoyRb?TfiPyk8eHKt!T#T6aU-5?lbqBd(^!e+wEZYulu;3>kc$y^@;zVz@A^8XP%p$ zqp{ry-hwt8t)A;HG-LIN|J`7(8(u%W=6LOi?H=$Bw63&zu6xmp)hGVd@qZBPb>8d!ezZx2?IG|`+C*AC*TZPW>XTfLfcK@1qm7}B zrp5Ls*j&rf>bV|6GghDYKMvlTHj*}+HjEbA6JT?#Myuy~63tkB;{OzQ5N#lB0BuiN zY)^w-ufDW;u4mAU)hGVXf_J6uLfeV9BQ3V)z#hl;ay^f3tUmF70lY138`@U1Eorg6 z2zGz9m+Q~y#_ALQm%y9RHl=M$>qm?2W$;F{_Hw;~Zmd4>e-+$^wmxk=+PbvZUIVW~ zYj5peM>kfV_`d<}Nn4Ay25og(Y<~gwptYClujt0=6aP2CE7Mk@t)Omew0$S@7I-<@ zvb6efd>hSJed7NP*!RIp(Jw(;ys*6sUX0e>IrJX7vHB#}-@uE|7N#voTYwhZ`{4O$ zU1{~YK0q^8pZNbBJP&Pd+MKjGXtDhRJUeYRT0Pf?XvXRj|Bt{k(`KU0K%1Tx+sEMP zXkBRaT%VvBt55tt1$Ux#q;;S*Xt8|;_Vcp#uJ@l~GghDYe*yOMJU`d-vp+urjP0M; z{9Mq_2kUixiOpDj;{O%c&n^A@($75o>@&81=?ib~UjH>VWA%yuzrlW9>*uz9*83qX zwr{ZcIk2Ax*X#Ngo3Z-D|2wdsQ~P=KJ2XG*j_rGFe*W#};PqVp!Dg&J@&5tr=ktC} z?`Qjd#vj{{*!rdOB zwmWSIZ4_++?LgWT+Htf~Y3I@|rd>t5fp!z^R@xo3yJ+{)9-uu;dyMuZ&3m8sEble1 z(7d+)mFBh7YoOOIuQ^^TJjXpZJ*PYu+{5lY_l(EdW8~VpCjNZDEcAZgj=ACY&vR=B zyGB0SYFhv;KliZ^*q>dv61)Bjw`}TupDp+2>-E>ZL<^VSXRm?Af6e@UpS@XNPx2!wANZ%;7piL682-boQ)}4-{$cOA zs#-RMdk->BU+o!tGqC5C=W2cI&9ND)PsZ*KUf`mItFgC$FSW|K)jDhmfBBu+s8m|sZw>bPka6uZ?`^Odt53$>7JSQvi&SH82Y+U^bE>hohyT`n=4$L6;C&d| zIDNGz*N)&{hwguGYS{^{)?O_;qnk^g#P0&$a!|LbmR;eSJaBeZ%Wm*@rkTE~Wq0_t z)MA{z+LLP!FrN=swd@I3Ywws0KsT2@iQfx6$B|QO83_Nd_gPgfgWx?b?Nrq=7``jD z7^koHLatFg!HOKWfJ3Fzk1C-D=(W4}7N8haA_qn|#h#@-L^J>I>kulB_558n0Lp4Hfs z;c7lJWPj#7IRM>U`Xv59Fy9SUpCcRu?>2bqIeRdCtMyN;YC8n}_aon_YC9Be4&(II zp4^9l*ZX*vsIcwB>o8SZ{4R}yN-m99D7_<%Te%qJHJ%bay0xfj*)Tt zYEQ0Xz@tApsH){yxLSKZPjDQ%x%5f=@!)6on0ie=0lvUWM_09+2>;K+k5;vu1V4sa zjMGG01VoO=H{1Ah9fM^v?)319S= zJF42wf}6uQeYGd|+2HLuUsl!jJGk1X?d+#>(9NY!;?D)&|IoY-gr5h0bc3lsTl0JP z2J2m2)p9=kWY)(m)p8Yl=J%$4 zXMQz&&%2MWYPkk(Eyn4qJ@=n$!IQE1o~(Y3Ux&?DeKPj-;9l>}T8(`J{M3hs_Ky8V z_^k(xug3l(+_8<*S9^SK0`uKP^Zh#-c@z66yw4GfR-gSo20!@go2p~z}K0;TgEK8K!>+&32ZIO1?<#}{->67>uz`Jem zYc;nQ;d3w4v#RCK@XbFxs;cEBxYu3d^wpkRFN5FwVZExBSKw;xy&rrP-CX)4{x$G< z_k35?@;dzMJ}XtVyaDHDPpapkzrcN;Z=Al`lk2bGr4Q|2<$4p&&x%&Z>@B!H>h)vx zHk!HhNxpZ$BO6~<`QC+Zu-6jR{&^2>E&6HmSvRqNgI|B*Zq>TH4_EV9ygv2^=;qQV z@qY&|_3($)*#CgfG|d9lx_k(Kdey$`W-osPKaggezS@)PWAKlM?p@XL30$pbJGFd@ zZZ3Ti{~36TxlgL@ji1B!zyIB;wlCm2eg4O)mVd&H(eHf5(4Jghf=5k0u&U)NxLSLk zMf?lhT>2#bYw(hPKeDRj-|$6$d9AAD8~AnWb*a|rTlgGDHrC5JeFry(ar$ac?(e}< zIvri*`VV}?)4Hvb*dO5fsMn9-k7(xBC;5H?f8FPWD&NoWNz?qeUhKcXU!A;6Rok!n z(~Q$sdwhQb_v&zRwQjySSKF+eW7q+kx%5f=G-^Zluhy+2{HBgiRJC=2&(-_Gsps%=`hIgHa+dvbRHfAh;(Rj%pavuw3ewQkeH^--^{+YD&t)+hO91i#d{zsffg z{OE&UtZJJX&K|02n+5KB7~}NSo?NqnZ}{{KV0~TYLpPT`iSG)2?XerG zTIPq})bXln>;>R{rly}ZTlVgPV0G>Fu@*u%R-gDU47M)oa}LfWwnf0s&w18!bwf8+ zpZIqNyWXz5$Hrq6+oE8Po5!!7YcX_V^@;!DVE2*x$vx;^jBN?9`_p||&$T3)vHHY+ zDX`~{=aJ`{=Ui+{V{4wjX!TsnU^7;q_%93gy5RM}Yl_#F*p|cQ^`?34YRxn6AD72w ztUmEy0qphA>!R0Auc5K6h|TM&*VlSoD`7KMpZKo~_B!wN-g|=ghS*lY=KaF^Mm^W6 z*o@UD{;R2x%lnx3I`4V0^}yzRu)WXVSI1_oKJi}z?0wk#vG;KA<*}`a&HHJ$HUz&=O$JmoW&&t9>ui_Pb;(`fZv z>tQoipZNC%`+VqgqR*B-W5%{VHlI7&%hd;)vHHY+1F+A%_Y-q3?VcjnhS=_g`y5@b zt1mWV^+~Rcz&_9aiJ04Iw-veiVY?OH-uE&aV>4Et5`+gT}#_E$?yMoPi3^7O1jx2KRhV2Nrp99tR z*Y4Pi)hD_30Q-5=VZ!uKFKv0?B|F35;K-IrpPq}+b(cFkF3veC^lpDNv>gFbL~USNZN=Z*KlmZ;qBei zMqo2mpX3?|HrEhh2GIr$P4m5?@6rB6`;pdxF=wF7PMeRm5Um?+G1`)}WoXOOR-&y+Tb;HR&3m8sEblem zBfPeIP4-&qHPCC9*Bq}Ep5vaIo>QI+?qT-BPlYD=8ESUeA&#r#|cqyb? z!C!&w2G=g;)M z)I45tf7gL8zfXQ2Y(7ry=6d|2#q;m#KZL8v&G!+Q{~GGn=e)iHFG(N%Biw!B`u+qq zznbws)4NyHtw-+f{!vT(93ET>3ZEOk2?K@C1D}owhpz%(v=#6AZv!vu836ZQVm*6- z{deeU)-#aavDK|dKB%Q>AKAjj$WMS%Zu5J1$9KHHqB*|yf41!Y-U0b`Y+U_aXOG`; zaMz;)t$exu)?v)?ZgyV!T9e%0H{r`;b~@N~XG%@~GvMY^*I(}MqfiU?_fg2re_jh` zPugFIro-}7XC8lFg_^$?z_tcwaB74$$Nmorr%7{+E$IFIEzZ?3{apn9{ucSVt(wn* zTh~^o{yu}V!D=>N*0&y+&7-Xk*{${WH0$?w7?{g?&u{7aJC3|3wOiu_;BD#6qwN4{rwC(($q7qu{*W&qg%QB9VXki^u+B3H_mml zUgP||ChCdX1MKf$P@mGuwI|qo+LCK8u(_6{CD%akUbN&I1ornes3+H8u=%w8%Dm)R z1NXAOhs4+1G}plEi@%pe?C+%T_pt?-RVMf12BuRVKYG+0}5jRBYAjfI;_JwE$_%kjp; z{XHM*8E*pE@r>7=@g{<`W&ceATZ4M??*}ex*dJ~U>hU=MT-I%C2IIA- zhJ(S{Qo|u&>rhYrL&0SYhr!LS9-qU(Wero{)~TKvjsRPO@!C_vkzj2{w$ACJz}BIj z{6~Y!8jgWmgL-_91(!7(2lrU0r-tLf)?mE$)Nle=TWUBFY#r*!e-gN?;bgcqsK@6N za9P8tEg$vNa2nVejMttTP6umC4QGI@LwzRpuGhpfX`B|E5xh61c`f(4qMjPg2A6gI z4(>Xr$LAbyS?9TMkC%GtJP&N0#%oWVzXxkeuJgg=co)FUr5>LP!R2@twS3ew-o;?Y zGhTbfy9BH)d-M-rYq*4#{Fj2u8ZLudgL-@}2bVQm0e4;0Q^S>DYcO7WYPbrlEj3&X zwhr~=zXn{^a4mc?O+7x>fy)}Mhg*YsYPbPx4aRFv4L5?dWqN!vD2j@J|?s;++y?SbR5M0*z5ZpbZ9-oK7Wu1?}^-)iq zkAkh!cbtwBBc zp97aQJP&u>)#LL5xUAttxHYJ!hChR?!FcVd;U%!P)bKLcI@FW@6>wR@t8nwH$LBS0 zS;Oma>r_t-Z-A}AcwWo#|!P-*8OknF! zPyU&~Weu~youhhuW(Aiu%m%jx_0%vs*cy!2o*L!=YfBAtf~`Y6=gC}P{%f8m+C5J` zg{vpeJm6ty;q$^hmuBK!+qCra!TGQGcXagCZXVy)X-nMv;4;qw@G{SWaQB5^*Q`fh?dI`)q_)H@3NG_31~2n04(GpSJ^E@lkMBRV ziGA#40DP0jTaC-){`>&f_=f*s%UGd`Qa)ic)SVB^$(AYby@f6KYz@$FA@T#uFh z<(ON*%Qe^%?p{m&t-$)IU)7r5)?jmLb6x$cLoGFK3-&mNZwEGK=CwVzzOFlRH!8U~_77UALiE%RGjF%VRYZ?zJuF@i4eP>c$V& zPD}gS-T2}3YR0>dM$tUh;iJLUnEJ zJ?HU0VD;o34=%@<058Xx2-inF<4giCMaww*f%Q|*IQxRtGtOjiInDua*DUin5Uh`S z#yJRVy&30Vuzu=G(pD4pNq2O|y!{Cl{b!%M@2dib=DPUvNGwu=Ka@-@~`l)B! zL%`}8_h_)|7JdxaW14Y}1?!`pagGDKPMP2FVExoH&QW0XRy_@V3(S~@u>Mc9{4WQq#oskhi~mf_OD%k6xH-2beirboG=0>!Y2}=) zrQ=q`&5l;)nFC(tnG>#$dh*N#R`=yxT>lwqKI3|P<^`)cp1#)TzL^iK%@~iNJoV2H zwvMA(wJr!(TY#p&$5bu;-N4ov-W{$_K676btmZR__pe3h7lS+Q4q)Sa9$p+x-P)I+ zUy`;6&HUzYY&HEd)>7bd?4{x5*vr7xJa!p-S-4~GT8zCMntH}w9_-lWH-}@Z>6fur z0AEU;Ynjt>^ee*EoTvM9Ww0;zr?!=7YVJ>Q@~i@O4ykWdxb-bdY}RHqH1*Wi18jZ9 z`o6~c)bvZ<)xl-nwczGmv*qKO_XMja@7iGV8oLI~ylVRS43YZnGtR-B)2>7KA#lgd z{y!A1k9zj)VPJLKs<^}9Wu7VUGS3lkebkfZNU*vuYxVl#&y_eI_y5sgHOJG}8r}cL zfVCOpn#fcCabU*{KOSE0sT1I8<$gX9Zk}>KpM<8KwLKZ^&v}^N9FDD~U&cBGT#kJz zyd3*9xLUcNPlr2pxu4HKQ_t9Ef*srZ=5TB^{WA7hVAntU`E0ma_VYPlU+!mZzoV(S zpT)^@F4#F_Kc5G;zH&eR9!)*)mB=HGfvg{bapr z`g=XR0<0$YT5~1X^$NcVzBoM3X;;II(NCK>owIskt_8c!;n%@k>%7mt9?sCs_u2Yt zkM9j&^Tqc@aD0u^?wGEtw#3~8F2}qXUXFPS+&u0NeYGd{R&Y7yZH2FK+8xtlr_DIm z?sjmwc6Y!XGkO06cmHJV?t~kopSG->dSdPdFJ9!h2kw5%dfW@odg!Y?vG;+?dE8(4 z8mHYc-E-O!_aJ!jV$6r&<(Ln{vv&GwPwXS$a?D2yU*ohpru$l3;vNUPp5ae`J*N3g zwU>c)E>sKx(jusOq@0Xwh6KMU4J-FVL%wfH{|HfQ(?V12UAFM=~?eYMB;&)_ol zC3qS8GCapfU+rFJ`yjpwUZ4II`s8~JtUd)F&%xKh=1|Y)Rd0Z;Nt^TWyjM$(zkt;; z&R@Zf<8?YdZ=$JZjotzqr=GaC!Nz$VPTV_a>iNv%U9hq0Gx1YJ*8CorFn(5z?@++L zm-78^#`+tYT)&L^;!VXwz4 z6?kQ^*Y8#7y?*=b@e%oQ{g!{x;@dDQ^M4FCUtZBaDRib#)&D8FW9GfzXK=N$xn}16 z9Nk#+JB~T^_xM_u_Kf#WuzvdcivDGxmvwxFUe@t1xLVm#$Jgko!*R^1zje4C+C4tj z{BN+eg?|IqC&%wwu#0!T>5*QwdFXwzt*DpTAk)`?m@ps3wv)~ zljgDQNuQtH`jH$tw)*^3aDN`xzcm}Xd&R#q8?L{9XEt2_ZUxuhzcU-V+`lth@;){0 z-<2)x{$1IU`*&qa?%$Oyxqnx-NM_u?JE=C~Qr`(8)5+B7tuGv&WvH!( z`C4}S)X^1ejvSx)!D=}^^0{f|GPgLf3xLbxvmo4DIX(-4kELe)v}NrU25U>6MZm^6 zw{G<1@#zj%(W)0O`7uU$S*{7x@!Pb=HvlLh@ z$48r*@mY7*DB~{!c7LSCW#MYx$NI2myq_-zU!UH!PTcZf_1ybb02`~GxRt=h`ELLc zw=!5g)l)}L zuyN{%>jgH>f1{ANb-?OVTIbWcU}M#-#p|P5{MQ3FHTJpZUvId_JL9eoR?{!XTP^V$ zfXlh`g8Ys=0hy=mVCt`__5Ejy=G)4m&AE%rTI zcDkm04}h9$;=O1uu-7YV8vvH~k|y?fmPuLt&8wmyBHWk!;( zyw~qtaPNWR3vRv%1@8ktq{W?^xpMC@PwqYYfF0M~d(UY4F*ILZBgeMv_BkHoz~;D_ z{pGX4zF@U!Xg(XrN72k>Zm~63Z)&jC)G)E-XKxMr(eF=74U=1T`_ym%*c_?hK(JbB zkWZqS%NoRqJqT=U_`zW3V~wtv$H;r3YkVkN%^3ZRao-&V)|MQHgFS}fQ{aw~d+HHj zpLz7tmOMv-wdGoM6xiJA@i`jo+GXDz16I>F>#FAQa^0M-eQG%tY%RGb9tT#-JyDyQ z@i`W*TgE>g?D)xdLK9ErC&FF7%=IL&bJb5<@|+CTmbsn+Hn)0wP6fMHGuP9=YWij$ zt68i2)VbQHw$s7ZlDVD%R?A$qsTrTW?>^1=XM!Ey^I2S;&zo_4eQikdeD)mc+rsXZ zjSAeag*`7fF7PHT?EB75Y1Y3vea`vcF=lzrUr_Mr!Iu}@+_%*D0|nRriGu6@YQe4f z{epi6{=DGYJF(eQkGxBb&sO7e*7$-ozDSKPUE|Bu_|^rt{+$c%{Prri_E9x{c)_jz zq=Gx&(+jTsni{{p#&56jJ8S%*8h^CLpR4f~Yy8`SJD*<*?tG`^VB~mcpQFZ?XmQUK z_mbD!K3u1r+c|LeT=`wdxoGP7*^%?WYQ}v}j`P93ysm5eJx$GhD|Q_B=x}^Ir!H*S z?7fy>On(W@m-auj?DjclF9n+;*YeB2YSYjhSAG%AT#hSF>=j^R!>?@dT)(e^dycA? zpNm}$*Os_zz~)KZwP54a%g@EGgKIO6L({k(yb*n#)o*~abj{Cg97BJ7atv++dk$N- zYbQ6xwQ-%CpMC23BiOofJ-Z34mg|`|wQ}uUC*%FR)+<~T^Yx$qR#;a#-?gYCBb1lCMtfp`Fr<%2APuXYOyTOiIet&ik zntJx@y};%VzI>eVG0cnlI~mv}L!?8a@U#NA~OE zV6|yzjw^qNW-iASC-zCOvEfg(c=p}XaQB^iRa@5b8L+m*JqtEZ;+_K==U&y8xaYyz z!e0RQ!!P^wMR@kBW9Y9>j=`V7-b1Y0wUZm;+PF^6&pvg%1h%g1*O$R+*{|BvjL+J; zPR6@mUjaK_^1TM$7~QeGcgeF~UkAHijnS5x-vHP5&0o;evtR!THeNkz^Cr05uW!NC z^v!-%v-a#M`;7ZG*m2A6Q{O>T&whOutd{-yH?S}Fjkfn_YObp|wf!Bu5q{zSfO{PC zdBcZrebh6@kH8<$v`uNWy%+qrrD=0c@6)SgAAJGN{_~h+ynmu=i{FR0I6 z68A5#V`RU54bFZur~dk64gU@PoR<9}HzwnL1I|7#$Nd&vTl~Ij`RP+$zrRPo7T9T*s{w@&k#!(7R+3fLUR zmwRJXbZyBYFLRhH^H~k-IQshjNuK;Yz~*<)l>26NbZyD82H5={z9!f;%D zG8@`*{qEV){a(_X=FvCTpS8hWe>^{}Pi{==>jkzxk8^pf*Fo2ode#NIXT#S6m-Y3A zr#@||Z~c~TedaWezNxPd*!mpX`sBu>z74?E=dme|&xYvQQcqv7`z(ATa9LkJc z-TKUF9(_~aR$%L!hGup|DSBwAlV6|L>cL1yT z{f6&HcLMwJT-CNCZ6BKFnb@2=)4TsXe|Blv?0vr9jed8UFYSA@?Dp2+=goV9&5`H( z0bn)f?6~q>Y36cVabgF8jSU~v;&~n(3|I5q%$g2?yQbQ*4nx7}@fil*i)Mc3DA&h1 zI3Me<&$z?E){${XfYlr~You1Lk@Ly;Bf%brJSXf8R`WRIxTzWMF>!p4gSJuhS^v>3 zo4xBlmVO+~m$Cb{?DnaBJlGsr{|R8VtiOB=&0LNvPV7XmvEh?iJnOh0+_hHET=xg7 z$7eFwwKkXYk?Ui<*6x`08RG!3V`i-n1gmAOwW%4O^|1DG?Oa3i`5x#Xu;Zt`gTWq) zJl7uzRy%|?h8XwkVPJLZJ)B;i_eoR0Yqj!f%RG+&*Wa%kiKd>hjshF6o_dZ3m+zB~ zfvf53_}bK*U*5~uXWV14Id1v$p~s=A=Y7)gV71(%PXzn&{L^*Lr=E$fEq-UU{OZp!XQOLNO}_&> zMy|Q%fV~!b%{8a~`eeP%1)omKev%uLanA!=U;Uo?dvtB_JHO>uzqeh0t}XRl2rlcp z2%h@PslPs{?_#j^nb-A`8(#h z=G0%G)ORJ=`tqJgZcOrB1vX!cHu zHT3e>?f{o_{S#a*bG;kv%eiX1i*_H)I>o8;9Vo#r)%xsv04 zusK|-@;rP1U0ZU<%N*v)d>#Zlj=s0k%VT>OT+a0oxLW4=IM|nS)%F%=VXowO8f*^dS?-Nz(6uFpyv$*)%;#CK>A}g(c56JGo5JKa{Yd% zrF;E0r+M_v_2*si=``nOeR5+`-+N%|bMED_{u{cs)bl>rJsbW3xUBE*@YJU*_5Guz zTc0`2qi^c_5Nv&pZGCcMQr|~l>(gEypO4YCrJhf~?z8Yu!DW4)!Bd~M)c1Kyw?1>4 zN8i-<1=#vLe%2>9CiVRjY<(Wv@_hReU0dq;3hdqp{};Hd?`wGK)0X=F-O{bkoaWIt z^?d`jKA%smPi{=c`3_t@lY9^FN3Ps&{sY%XeM;*&mfJ?JN} zTHYu90`}$ktL|NTRWc*pd9)}#u+2Cq6AmHa2{L7SB2^2zRa3 zGuMT{>hW0^>{^@4`N;LLUTb$u`;4&&*fF!#-N0&DYi(-fdRV(WYiDmh-zRm)=J=^^ zQLyWs=laFqYKNfZebVA^b?f!Ja(Uh-EeT$$l~-HlxfHnmer0Ji^^CO)*m(8Svn;rL zpR^oYO<%{?rsn)kq(;}vKJm+A)5o7h5$9)7{5iaBX}-3i`Lifn&~M$sThebsGv9Xf z{%naqE3gtd_^;xt*LbfQ?^om7)cEc-KBC6Q)cC_Rk7#z2DUMyf`5>*I)bMHNI?(uTGw@PZpZy5Rbct?_XM*MCyM_20k7Cl_4*Lkq6|;Wa*` z#*ZtwzpMJhf*XHk!9D)x7F_!!1=oI6!L?ss<2Tj#odq}k!GatAWWlw+P;l+9)c6}U z{!Wd5Sa5&$_Ll{BKmJy5$Llct)cqyzSa8ShQsXn#_*@0=ga3jBH{bF#-m~E5?^WaL z7TkP&3vRxCHNHu~&9_;N_b<5f-?HHTF7oYbe6ND*Kdj)6zjwh|%EtJDYoFNS-UmEi zSK;LK8k?V=ni;$*n!3M}LB1MzWt#fITz7p=>H${wn(sAOo}6of&6%HZF~?eH>dDy? zyar7@IoAfOo73yK-1~;(tpj$Pv2boDjdj6lQ(DiK>w(q$e8W2YHx28<*QfXQ`{(}C zM~#+y+y-D{)$_cwAz0m9d7so5PAdMz8{ct&bMM^nl}NfrRGh+ zYGuuv!>xHU`qbPXte%>;02`~Gnzsb2mo;w%PtDp>^VVQ(*6jUIo|?A>yUux@*$%9B zX{V|4*dDH)dF%jIE9bEzyuS83p{ZvcJA;i?&pdVktC#cG74AIPx()4_$8KP4&cpk- zJoDHC?Am7@1Hfu~(sDfZ0;^{p1Ho$LJO;t*$73*>dgd_%Y^-|bF%+y`&SMxn^U$7o z3O9y3fq<8x2>_bLj}MTJnwsJ9hXu zuzMr%`-1gR&p6}3>g6~S;OZI2=Owk|od|a9{Jxt>U^V^YvtP?6EU z!H$!h2Y}VeoCmgilJg+An!e^VUM)Ef0Xt6kp!WV`!SrhJp8_^M>vsfL z&9zDmM{1|#Gm68(YRPdl*fEpi7_eHI<5;+QavTL#OOE5g<_$jqtWV}~B3K`F=W!gp zn*Lt5PX=#8pX>H1a5aA(O#hR>YCi9uO7HWzz3=xehY>y_)!b`r)jTd-wscd&=HD{t*4cG+){uY1!>F??=Jr z$Y&alfz>=`-LLWoY34Gw*l{1HAKuFQM9XIHxKGhPP4nfr&$R6J8TVPRIWq2ZV6}`Z zf0AY{#}ymms;FC@(S4FncS~}&7FJcYhbmk`Rm~Net82;J!}3KuyN|ydw&Hx zmNv(7>^H$4kBt2m*!9fVZ-dn`_B-JE@qQOgJ!8KIHcmZb{|)R|+WdV|@6&%kvnFjj zH`;y=-rvElkG5&3&+ivJUeCj|rOr>l)*1dO*fo>SMcn6LW47X{lCJRu|EFg2McQV5 zbAN`Wy(e|KCjSJ_NUxsHjK2hbOmkevGEPlj_lWz&<7}Urz5-iQ?n(awtL6T!P0jf1 zLHCQ?eQ!Sd+~DgbsEliHvL>vAMUlekNll*9pJyfjW@4cpX8keY~Iw-5w7O@ zo{ZHAp0V;g;#i&0wPmbn!H#7P$CB%lvATfEv8IEoxh@%Ndbndb&Tv}hH3Pb~j5Qd_h|nf*?SygvCY3nmfZ7!%jcV}aJAgq{d`N! z-1F00o4wcI1?d-}`O>~{%Wj|ii-64`_j$b=_&93IJ+M1mEzdWLg4JBVJl`w^H;*>& z7xD#Yj&D8UjJE{19B)bZ;l+4M!PWA7vozQk^_+9dfYq}Gi-XlN#&Y0tjOF3w7%RZl z@_e%**cf$>m!Ao$CEvZy<1N+pn z8`xSL+d1tHR`Wd0@li8A$H%plyC=+NpL~0OU8CgN6Rehe+SH8CK66hP@0t$=drrt_ zgAWBeCu%&N(&h3%4fy_0yi3#)GZNvCJdaKQ&DNmo-gOzSML8+?w>)PkU-Q5Nu73WgfZysp%kaS<}IAHEYV+9RgM_ z*X}TQ>O2%IH~-=EsdGvTr_Lkb)~UaK+EeF|VC!@&^T_pge=S7s{?&!)w|3U_WzIcU<4QUjX*~P;Y86 zkA4@TsptLBMPOsq^BLyFV0H6Z%ei2+)ovfF2k?f{!3_uoH()%@8c@B8vw zXy!7vII(wuS1e!ZdUBi$I#SsEFK3NtDfWe1Xw-$^FgqhwYbNh z0;^?TJPppi(3YB?0rxBR=d*CNjPo4WSoQ4B=fUc(!ISjzjQb+kal`-I;;G{$ca(})KS2z9@dU@vg7qGSF_`V5N`ztLzZ-G7EvQORy z>!)tcH|W*G@6cz@zYBI>+PmlfM*lv|mwWz$mfb#M{T*zM?D>Cy)$;k2{5_hv%q@1@ z59u@RM=hJZ<95^G|2+Dy@9RX+MRvbls)} zYtPTP{YDP8%walkdG1UPSKFplm)F=C;N`h9BV0dq>+xq5)WrUrLiX^?VAssvJv=M@ zY&2i4$?Pq=ed?J5Y>w>VIl*dvH{|_7J`2rU<`z4yKM#>{=Wf~T9d};(`DhupYs+q* zapwn{BjYXrR?E2Zd1&TxTybI-0$X$V!eDcSF9LRr!n=X>&$-hbtmc|$FD(i;PFvy@ z1DiKKi-Yx3_uTQby(PfD57cIy`)WzBdfo>v1@`AftlL^$KQ;Y4F0Pexw9nW}gUhj( zfvaV&Yg02m>+f30J?7@KPrhZru0iIq99S*+w5b`NedIBBp8mYY3Sjqgo;jW8ifHP& z@2muN->WBXWw3F1hA?gwH1))-3N}_fajSui^RuI@YY#N_#H|iCRy}cRfQ>8v{nVOh z>WNzmY^=J+VtINsu|K!t=kSiXHe;_1_t<3}*8yKoetrF!9(io*fxTC1^XGfyUNf9? zZ*cEcKI61!Uh9LcA@}1xVExtOvjN!Jrf@vWw;@E+3_F*w(~ zeqecQn}T!Q+XO7v=5=XvaOS=lSRUIJ;LN>0Sgy@$;#OdfiM4DAmiOeZ+dGe~!JE-r zqqR7uxy_+(>e>ce&VO6Dn(>*xTKsnapUaq8uN~oPKI6Gpd_S`j+~>s`Ie!zkGn#si z#V%lD)#I}(_`xFgZg90S_wI0Wm){lcfu^3^dxDKs&-X+Fz>cZSd2LUx=GY#Cf#5S4 zJI}>~;A*iC2H#oOhrrd6e<;`(b@T5{-UOnG) zO#u5nm%8@7>D6*fCxP8xIX3%&)x4g?XMebHxsFT*>!)tciS%mmKLD(jxf}?tugO7Z z>Qh?h^}%2@V-t4>IB||6*DvE73NFVvtmR*>>)~kX8D|Pu&De}{1lVzma~!#T8RtlF zInGgV{mWx>G+aI790OJ}Hsc%%PMqV&^>du#=uf1b(BkgFlfdp%-!u8QsZIu~>F@RB x6tMXQ!;RIh=DqAT&IRv<_WFAcor>)ycGQW${vQiLV{HHc literal 0 HcmV?d00001 diff --git a/src/Ryujinx.Graphics.Vulkan/Effects/Shaders/NisScalingHdr.glsl b/src/Ryujinx.Graphics.Vulkan/Effects/Shaders/NisScalingHdr.glsl new file mode 100644 index 000000000..b1409e32e --- /dev/null +++ b/src/Ryujinx.Graphics.Vulkan/Effects/Shaders/NisScalingHdr.glsl @@ -0,0 +1,125 @@ +// The MIT License (MIT) +// +// Copyright (c) 2022 NVIDIA CORPORATION & AFFILIATES. 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. +// +// NVIDIA Image Scaling (NIS) v1.0.3 - HDR variant for Ryujinx. Same as NisScaling.glsl +// but outputs to an scRGB (rgba16f) swapchain: NIS runs in SDR space, then the result is +// inverse-tone-mapped to scRGB via NIS_OUTPUT_TRANSFORM. The tonemap below matches the +// fork's existing HDR path (FsrSharpeningHdr) exactly. + +#version 450 core +#extension GL_GOOGLE_include_directive : require + +#define NIS_GLSL 1 +#define NIS_SCALER 1 +#define NIS_THREAD_GROUP_SIZE 256 +#define NIS_BLOCK_WIDTH 32 +#define NIS_BLOCK_HEIGHT 24 + +layout(local_size_x = NIS_THREAD_GROUP_SIZE, local_size_y = 1, local_size_z = 1) in; + +layout(binding = 2) uniform cb +{ + float kDetectRatio; + float kDetectThres; + float kMinContrastRatio; + float kRatioNorm; + + float kContrastBoost; + float kEps; + float kSharpStartY; + float kSharpScaleY; + + float kSharpStrengthMin; + float kSharpStrengthScale; + float kSharpLimitMin; + float kSharpLimitScale; + + float kScaleX; + float kScaleY; + + float kDstNormX; + float kDstNormY; + float kSrcNormX; + float kSrcNormY; + + uint kInputViewportOriginX; + uint kInputViewportOriginY; + uint kInputViewportWidth; + uint kInputViewportHeight; + + uint kOutputViewportOriginX; + uint kOutputViewportOriginY; + uint kOutputViewportWidth; + uint kOutputViewportHeight; + + float reserved0; + float reserved1; +}; + +layout(binding = 1, set = 2) uniform sampler2D in_texture; +layout(rgba16f, binding = 0, set = 3) uniform image2D out_texture; + +// HDR tone-map parameters (filled host-side; same values the other filters receive). +layout(binding = 3, std140) uniform hdrParams +{ + float hdrPaperWhite; + float hdrPeak; + float hdrCurve; + float hdrGamma; + float hdrBlend; + float hdrWhiten; +}; + +#include "nis_coef.glsl" + +// SDR -> scRGB inverse tone map. Ported verbatim from the fork's FsrSharpeningHdr path. +vec3 sdrToLinear(vec3 c, float g) +{ + return pow(max(c, vec3(0.0)), vec3(g)); +} + +vec3 sdrToHdr(vec3 srgb) +{ + vec3 lin = sdrToLinear(srgb, hdrGamma); + float luma = dot(lin, vec3(0.2126, 0.7152, 0.0722)); + float maxc = max(max(lin.x, lin.y), lin.z); + float hl = mix(luma, maxc, hdrBlend); + float hlAmount = pow(clamp(hl, 0.0, 1.0), hdrCurve); + float boost = mix(1.0, hdrPeak / hdrPaperWhite, hlAmount); + vec3 outc = (lin * hdrPaperWhite) * boost; + float t = hdrWhiten * hlAmount; + float lumaOut = dot(outc, vec3(0.2126, 0.7152, 0.0722)); + return mix(outc, vec3(lumaOut), vec3(t)); +} + +vec4 NisStore(vec4 c) +{ + return vec4(sdrToHdr(c.rgb), c.a); +} + +#define NIS_OUTPUT_TRANSFORM(c) NisStore(c) + +#include "NisScaler.h" + +void main() +{ + NVScaler(uvec2(gl_WorkGroupID.xy), gl_LocalInvocationID.x); +} diff --git a/src/Ryujinx.Graphics.Vulkan/Effects/Shaders/NisScalingHdr.spv b/src/Ryujinx.Graphics.Vulkan/Effects/Shaders/NisScalingHdr.spv new file mode 100644 index 0000000000000000000000000000000000000000..e9567dedb304466e4126ee54c0e6703719b7f51d GIT binary patch literal 59420 zcmbXL1(;n$+Jy}t(h=NUlVHIuXn>I5n&3`ICmj-qO$4{#GB^w{xVyXiFu2R$HaLU3 zefNEOubs_)`~R=+`p%@L)>@C$Q?+Z$Ij8Av>9ks`6ou_UzUo$m2p)}CC zHv9EA*>aPm$4?l#bng|H)nWQZOVdxE85-RhUFbX0MhzY@TEkCMH5$LtW+YyHCPkd} zx3%!UPV^5#Iya_ntiQ>i^#=@Cf0K0w4H-Lj(D(`ajBFb;a>VGiK|@E3Ya23Q#GY;X zOW zK4Y8yEsaI!M-+Hb`rVs&%EprPn~WJebiENHC$x?0KX~*W zy;lOSJ$U4hO}5&Q0bP>nzsWJn;5W7z)6$@8j9sq19=gWZ<=Zh`W9$m;n65FlcRQwQ zj9pROr9syiJACl?wm~CC4{h5UWtv7`H0HhGXsR8X*|Uy2F26JXih}U*xUNp$#sbm7 zyEfLKZ_hVu+~6UD_Kv>hU-f<3^=^&z0PF5Kc;vt_BllU!C0?nyzT|3YtYggHBf(u7 z>z45&Mo$2{R-4kdujyuE#*G?0vOT=W(7sg7JR2F?jOo_cSlpcN-h-MIwx`3dTVs>b zwolDAZDU(({e})*XT-Sn5ywv+to;Bn9{-xV&Du9N7tJ-s) zTb26%t#RwpR+bP9+tVimA8+($%s)SCNzUH_i6rWYY%MZZl8zu!#$encI0c?=7^wnk&c!r&&Xn#s`zN z=7q+5pH0@97bCT=Cu_|MkoEtG*8a@XxzUvqw|Q@E=Iz{=3chD~kKtKPUW|DFCvo$M zV{q$8Zv)Y zxoP3onUjBf+ej`$W5%_|wlt<}*T~YjF&%jDpi#|pSUyWT_j1~_&j)X<*cWW)rG2rc z{}8TL?fy$O?crS-%fV%X#og25;_?h;97-tW*V^&+db5v?R89sFUg#JFp zde3oP&ZubKW4v~pi#C4TkfvYrE|0VN6?L5e^_wgDx}%4hx~bpySKahY{R#BZ6Gv?| zVr1LE5qr0d9AD&ruC5pP->&P;Jlz`ax9i2btuhs3bZdOnfvtUh-5Q^DU~8|frST=U z_BC#4{01Jok2UlDo(Gt*dy6|ax`BBX88NJR?P_UE1)tEgwKQfE7co7+?dM)gV=-fz zb+X58quag*QAoQIw{Tn5e@ zt^>E%*ro9Zc>M6no@z-no-5P(t#y_a> z&;G(&8vm1z@p|96@dKFih+W&hAAW^%z7(8KgQm(_w=|{(4{NS_=f(_8?A5IOIL`ua zU(e2sdBEiKFgC|n5Khj5w=|Z3=U6Te9@adrof|7RadLKT^fAsS^Wt!LB;CLL&_^_L zw=^~c=bo?y*uBs`ZcC#T-u`~Lb7LsDdJKpCg?DZ24sVX-ncTi*nq!SYA5o087dZRp zAaM4~3E);fLfN)>GPTCD(Z-D3zB&JM^{wom^K1OF8o#2(ud4BDYW%tyzoEu&gAW-q zW?Zhv55YM;YZp>;o{yo8VM5ma6qrNOto-1n2BRIT*ur~Ha~|%-((C=BpZ(2la4ouqqVk={o7yibZ-1! z;pWWZn_6f3_!*-gTfNj;3eERd@iRuh87mr(4oz)V`t_Ns z&t5%f#)tO=`?5ygd#1*`CpM=q$5E@-m>B(1qncx=rAD=SjcU~zGX}TwX6&5w(YPfv zwRtO=<2gq6ipH`xwS_9$)P>foqD@n1OI9>%HQ%xo&G)Nny(^mUSJhUjXuelf^E1ZO z?6FW=tD^ZHRc)P$=KE8%4Jw-NN7Xi|XdXwkffdd7oN8NDG~Z*Yd0|W~zNb{%siL{Q zYP(jnxe9G)Me{wPe#0x8?+?{R(%ehFH|zm+|AbH8Tb-NtTkW_t?cRU2$KP{PyD#JA z+FJ_lcL5$VYtioa0&?q+w{Vc<);|kb&D(=$8I z(}P{=9@3_twjTI+%%hppV`fhGxj8&u+ST>Ytt@2gI0|{3+^lSE5P0_GTuF4=Q%4`P9<z&{|1X;B@&B(Sr*&M$xX#=0J%)0xpK6(FC&ukRytEJzJ*tEMY=C_ua;9pF)W}onx;aAKu@b2(g;LonJ-Rj9dE8H0U%&$GZ zvw@F(ewVxAJ3Cyh6~DC9G6%Z3^hx}j;67uYE^C=f`>#)wwag8_@tlLo{W%YO=+1AH zwap7RhjIF9Pwx4^M?e2oS=;<@wbl-5TL9f$`Xs(P_?@L@Ts^!8{ERnexhK3Q{NlF{ zENfd3{?V_CaqhJ5#f9MJFiv0XiCq}H^#i@ic`gDUGjq#oiCq+~k9uoSZ`QFFnz{8! zzQw>h-q5$Kb#eHlMf#VuE}^~uG<}glQ?a93~c>c?V-<{f)fva^? z+p_59(kJoDfr&3`TOR()7kiYotpNXZrvu7$>uL{27w|fTlI{cA8ZY*nC18xrE^wpl+ zYl631_2%*z_l1v{`OxwhuLajfy|s3X`=OazpXBQg?l$O*vevcX_f3CxS?fCRUk=>2 ztaV-ZcfY?_*18_t9LDLZJ-OEhKfBJWWvv^)FFotwvepgZ`lz?oYTXFU-1;Qn#^7_m zysWHs6Zn3!UR&0>Dg3QY`;@hA2JiX%4`rw}yW)-MwY4+rWGNzGqqMw(w04n{LgFw;kLZ#_6j) zxwi*%4JmW&0N-N6`N~>%gzKZ;TB~&@G;`~dd^>}`eC(mJ)?MI(x;~cPsep6?)#2c@2T9&17GzZ77<#^htah_=;JcEo&PFKkB>J z%i4B>Px02cvbN#y3w~OztZf9`9LDLZJ-K%WAMor7Wv)Ho+;_@#90}J)y|q^BC^U2H zlYFDW?Dw+PG4P*f_@vA~7QR{Q*s|8Y!B0JM-LlqkaB~=^ulD2~4?g6*^~+on;O{JT zVp;1%xIXHwwOaQ?Gq*m;w-@+@yWcNs-5Y+>r{9+O_kkaF?x?cXec=Ot+pes2Ke#!J z(^q?P?+-reyB*712f&v;@%*yZ1L69px7KPs2+iF3B;Uc{U%vRPtaTFn^`HMJ^B)2~ z^#0w;S`US!aRUtMyDYbL*3Q{{(YAE^9ptKKrlJl(n7>KmE^L%UaKY zAJF%-vet9q<}glQ?a6(fc+S(xTF-~y@msg@`R@X_KI*NtS}#O1w?4^t5qQ@cyUbkHdJX*Z`?oJ^y%zqkV+8 zKg(KggzKZ;TC4RYG;`~dd^dxi`Ei!Awp-x;dSmOdwp-x`^u495?KZfz8K>_Mp_yBs$wLJjm+FsW7AY32y)>>^3p_yBs6T{Q~#^1?J@X|6L&3ZdmOHhdTXt=C(z8TPx3tpo@@LZWo=Kv zXPa-svbLw;7q0tqS=%#k&kf`B)t+3>f_WAwYkLmx~A_$jNsTh{V2{3M#S=&L=kuYix7{Xb`nx8k)KF zNxs*?Jr?+`yxzS5f28|bWo>W5&s+WTvbMM2##o!a+LP;TaNAt(mG_Bv;A*SVT5Gku zi)Jo;68|1}$~F3wW4{l-YRw_U~t39#*1;6ysr{%hQ2!aRUtK|nYbL*3QKZ4h~`0aAP`~-jcsL9v$pW(xME?L(43w)`yr!8y!6>bjW z^wpl+zkwI-Iz@TT{tkbn`}Jk5f57!oZ>`n(Cz`qSNj^VRJAD5c%Y0M7KbmjRvbIj} zOYWJetgQt;nt0>%)t+3P!IL_7xi`nZ3tVk0{99|abVWCpK8f!J9@K66vX&{~yT8?= ztYs?rNAt~5)-pA`joilRt3A1<0rT8Y=9(71;|=$f>ogr)AN6`|)1#SNpX8eXe8_u` zmdA8P_}AY}UFM$&-s`}5%i3m!8>62#Th?WkLf2j&dsZ}K^@;y%VC%GA=jdExn;q=@ zU59$EIna&OC;oGSJsutxkDbRbwzJ$HY!0t2mn|su~8r!_s+`sPQ z*4nY351X<2#D9LU=a=W1=ceaqYztt!8SeSpTFccPo3Z-DzX#ashSv|TIbM5W>xs?l zk=LcxTCN4L8LLnH7Xo`d^t$*y&1-6G3uE&->-Dy^mTM7g#_ALQMZsR@z219Ic!L&O zFKpf~yl=GDaxI3UzYFVbRL0-L$KKepC#Es4!oed50q_&I#NKR-?L z-X7c1*v$1$T5B!WGT4mOC;rQVAH&z@5T9Khp~bcwHrLDNo7P&c<*^y7PyAN^`J$H!z<1z(JI!a&TWGPZjLrSkQLa_68LLnHR|VgI|Mj$M zY1hzVTMgUQ@Q!lz!Dg&J@n0Q$IsTW?E}>mai){^T7r{GP`!%r{t55vw7=7i zrNy=$wqxKOokQzmGghDYZvZ|5|HEmA(hi};wjs7j@Q%*yjj$Q3Py9CqAAtYaP|F+;A@ZFxa zEo~cGY}@Gz--^~+%e6f=WA%yu4&cr49Z1`Zwka*P9kFc!-){yI@-z-k;W5%QXm_vHHY+S8!i^*QBjZ>qCoeFt*j;Uc*~!xmvLqt55ug zfEU4cC0cLV3bfdUVp|@*9IdsMs|}m6`ow=2*j!7~mZU9VUwrSh8@9#ai_uzZxrSpi zR-gEf0Gn%J+Jdy6wAgmX)&t(1)>_N82R38%iT_Bje)Hi!4{dH*Y@@Kv1^3)qex z%~*ZnKL)JdZ1~SY^Sj5`#$uZpt+|(Lx&DUDSbdgktlj+1bR2j_{O6!8KXWf=1>bni0_E7Z!JnA% z^m6Rm;eYm+t{nRg_~qm`PG9ZGbtm|bw!Q97EqB4yI;!PvbaUyG_s8uC7YdCj{O4s zfWF_Cp9Q=Kf3-L!`f5+?OW>*InZF$SWw=^LW50rKE`1XJDtOFS$CYFM2maBopOj<2 z2KS!r-qcrnVqXXE@cpXg*l)nqd>d;5$JnEzU%33~yt8LK1I(?3AE`1XJ1^CIGCtuV5 z2cKt|!^>K}g#Yy5!(}aB!M`B4ar$acuCKwHetUdb%QtYfj%xW9-CX)4{yQ+&fby7q z5C8oB$@i}x;3v&+NLkyD@SeZlTGsXx+#JT~t3A1Y25;HrqO!JM;A)$9uui|Cn@gX> z{|3JIfjRCA{~i8t|H(fq^9Q{DnwOWg{0aY#+{WpvJ+a=%XZ+;0vX&{lvD18pZ>{a^ zPUz;+C-E)d?|z@Ptfe#jfWDJ|)}RY~;BRM?wRF{=+{WpvJ-NDp`%LjvS<94gwT_GzQ<~=Iu&r}clFC-HNDZ@6#TvX(jF5BHz^edb*7D-YSctYvPv@1u;5@;c64` zuh%jky1Dd8{QTe{2TfJhvH*O{q^-(z=?)*!SfIt5!^X#SXf^apT zrRudTgl;Z<62CBb!2DgyS{8vnx%0+-bIccod(P^o&Cfy-+Y8?Bke=mdxr@Q~|K|Gg z7%dLhN4;Ll5@_buC;65HuYOObvX-Ub{SN70uFKN!r{{XEJQtRMTbptEYEQ0Z!RH^n zY`HGW!PSoFU|p6+HPLTE5bMa^sus)mEc}?jnh|qa;*%0 z{pU5yT2_Ipb@cvjRdjRdllaxZXWsQgSxX=IH~p3=Ygrx6&ytkSLu^+PkaKFQY~JfiVsnQv`)|6La@_s=?VYSB-d&$@|S z7k>4zJC^IR9$d|5@%q^7qnk^g#BTsz^uZ6yu{VTI+i9M1T{eP0zWh3CWG`*%wEt{mWi8vnd;b1%S<80ttJdsRuG9AL znGbEOnRVI$ZVu!0)t=ltf+ux8yv(%|eCZRrub$YQ;rgi8kKrz8=GG_q27$ln_e_~@ zSNOzEzpNSiVE9XWFILvp3b!`n^wl2UA>ck!99OQ}P`KI+9UQ|pbaUyG_+en)ZE-ZI}<_~H9MTh{hBID4q9Z5-V9FvjVtJ-Nn%uleQda@{7t)jIlI zY$Ce3^hx}l;JG^ATGp}`{GuuDE^FBv&d({Awd@1;yAI>@)t+4Yf~Q&UlCqZl;A*}X ztgp-d=;qQV@dtoke&m|6mIL9}w_H|^eGq&Wa_Ohdmc4s0SY3O4tV!s`>J$G%z}97b z&cV6Fb|~2SInR2o!_bY@C;o?nU2oUjW8*Q3?Fg{P&Er?kbtJm6`o#Yzu=~jUSbgGuEZFnM^T>0}b1t^OgFO#DAM3e}Lo-&N_#Y4Uy5RM}Yl_#F z*iHa@z3C{|iRi}a6aSOIUJt!4Zcg(W8r#Y6fwTd%dR?cW8LLnHPX&9O--O~qwz^h?v0s@N_N z(|j(g=eiKx*rreM8T>`?#b~`~K3gtApK&e*FHGww*Cpu2>J$GV!H|4 zLhC5k&FIGJlU%oeeSh4b_r0|5sbjkp&G*?I<+=@xu0? zY<~XZ=RmEs_r~{QGghDYKLGagC_k6-vnxNtitRycey-KgHRmC0#_ALQhr#CZ^EW@E z^Rv3x9>I1i+|Tjqb^QyQvHHaSQSc)8`Z=PXHTs!jZ2!h~58Tfq>$x7oW~@H(e;jNs zKM(aYQ$IV6?Fnptp4!np?MZCL>J$H`z~=IEUOx-=GvU~t#^&e59X%sFgUwid;{Pnz zTz-CiAMM^suII2hM?VK|tsTebu^FpRa=if7?@nTFr`=Y`^&&Rc-upm3*Gt%p)hD@L z2J7ed4L1^ZLnYTM*sh0n^xo)IY{u%7T>k+(N540@nz*Yfxn9F|CA_202wul#tUk&0 z2H3hTBjys?#g$xdV!H_LbwnRqKKFeKo3Z-j`^C4x7vg^f?K;{mw7Y2!(H^HgM|*|# zCe6=g{Y=!)BK^#-6R}g#W}wYZn~$~-Z3)`)v{h;SXzSB9qis#wh1NzJO`AyDpEikh z80|>fF|^}oC(=%#{eyNU?QGh4vfwbW~%*DkL)UMoDuJvTk4JQv)< z?mhR6$J%4$+PWtBe(*Ce|C-OPJ?XphAMSRq@HV*L7k0;MB-rmC`&aBI6`Ie!#-9oI zzt0Xo2kw72AATv^|NcMxO1S^czT9!v^Z&S%$o=nWH-S5*-*;{b)~sedn-w|1EvNS1I_0@C|5k z^KAsC{D!*qxt`mD7oiW|89spxV7^_z=2tU*5WRat-FoD^7Fyyb!YwF#Z}ZhtepTF-;<37luv^AOnoPF&4;9;UY*b?cG)yOq?mKUrX79{m@vn#~u(H`XMxd9?K-yS4sFvwnZ) zfVr&O6I@+?$C0l}?XKe#V1KW$d9?jb=2d9U!FoIwec{$vnRpIO^ZFq+mFKQ=^LM>C zo-u=IUSp<($%oRcw--&@5SqRo+cugx2Ez0mR%q(&w)Pym!Obx@&3yX(!MN(xdOi1j z`8!8auh$1_^!LZO-p*lnn&X9!sc`4@w;K0%T*hDf0X2S9jh|5C{%+yKJN^X~ZobQF z{Mv$NO~=BWhsV%b_uIACEkL#hH@%$ZI zYNs-W$7z3>n!n>hf7e6J8m#{yn*P~$#vDvjkIx}seJ0UT=b>PA_lhxx(LBz*Xb01@ zC;o7-@m^mOe*{=P@kfFaZ!YcD>+#Z-xTC?&A^aGyK3UVhgI!a9H%EMq1FKKV4zkwc zX^yot&3xw2$MM_`+7o*sSS_(9gNTcXvuLJSUq$72iP2*6Xw(I z*zP@T$#o{!TxZad>z`ot_?!*a$KNFqpL4+KM-_YET(EKK@i`Cdc$wGv;H-)BF-9Nj z@*FsyW~}@0LYl`jHCzO)*KjeK`tj6Y%q2AA)KkNyU~|XkGO&K?j(Y*UTE@8&>^N7@ zlH)3{did2~=XH27_iMoVsGHOCRV{P54(vENU#e)B{2JcN%Puyc*<0jEE&&R>)iF*RP4^4eiG506I=F^s3Pl3%u__a$L9@jb-Xv>XVBC$-dkYDGhTbfdmF4R`|lmF zHK-^5yWnaK@4>A>JwES)t2KN8KZT~A8vYBm2IIA-h7ZBoQo~2!cWCO#|1r2)!zXa_ ztH<8_6bOFcf_z}4}5|E-UD#+wT4c*bkbcvFM5Wsgn+wg&a&pB7xL zVLG_^)#EcgxLU&uaMwjWHOvUM2IIA-hMBrhYrS-{mAW`(;y)Z;T7xLU*P zaBEOc4Re64!FcVdVNS5N?2oy?)}fyKbAzij%mcRu_4v#SuGTPL;iH}!<_BAY@!C_v z0$^>ap*z?*)RVslxLQL`xN}sG&w}7;4GY1oQ#~~-47LX2wWo$f)M=?6SEDo;LxdhxjqaL3n!PPpKDty#a=h9&7G+uk^Tn4Nyxt0Z2$6F5W zx~Ruzd2n^S72wvSp7DBv9nW~}8E-|fw(QZBz}BFi{40a2HLL=6-PPl>D!5w1YH({% zPYr#*)?mE$)UY~OTWVMXY#r*!zb3d^LtnW0)#I}kxLQL$xOJ+hhW=n{FkX9VSR1S@ z`(qt#H1*_P7hJ7jJ-Ek0JwEG$t2Jx@w+8jpup!tQjMttTHUeu)4I6{4Lp}L70at6- z6mEX?_-qEQ)-a&(QBMs6!Pa2B_SCRBSX*k?0&E@XIZw6(^RIcHX!kr>j9xu8Yz?l~ zxeeU4QjgEJ;A)-Q!9AwxsdIaa6U_KdeHSX=h!V6Zi)Cx0urTEh^yHK@mDD7acf8{BnKPYuJs)?mE$)UX>^TWS~% zwhr~=9|5k`ushs2s>f#!aJ7b!aBEOc4Wq!;V7&I!FdD2a`(q5)I@FVYEVx?3-{97u z9-ndGY7OJzu9bRfm;km0hk4u@Cs90BKFb1wR7H_xH;+LGreu=l+1qv6(^dX53}uUU`2+Rbw$ zy|%>t9bC_2{eJJjc>&6Q4leijVlj0w@1TaMvLGWUw{l+&BfSkNVBL zM>v)KG@80Oe4D7bpY`X{(L_E|N4HoRKPIdFZ{%g+G8 z=Fyh5JrB&k=9-)fI0>9}(3bJf2fH@edlwYG>RBtb_+JdJ)_V!u)mMTS=W2O#;K>4d%@<`=DOZb@5^=7b{|d6brom)2f)^o z@gD@&*YzPZ^^El}*f@39HF@o`u8)8n*R|HaI_AIN)irn&?%qoNe}nZ=&$>PaHm5e% z^>%u-%;O1g^;kU#_u7{8_$jzP>c&4#ua@{{z}1?bg?lYZ{Bv-9)Qx|dUd?#-(F-(> zb@+>5YfSwwf%Q?(m@k7*pj}bC=Y9pOpL)*Y=fUd9`yX(1oY&yhabAb(qn>fz0G~+9 zIB$aWQ_ncBg4Hw5+u-Ur@4#KN%;#OOKI$3gJ+Sp=ocF={spp#f7Fa#w{1;pu=R>&T zTwbi}M_{##`!U!U^^E%oxH|5qaQ)OX?gwD?jQcs*bqoIj>@m$a{{!ozo^ieeyH1(k zS780rGtOsV_2m5qTpj0IxZ`Ag-+|RK?)PA0)HCi6;Oe+P!u3>7ms0=CxhU%}?c=k>pV&0(Cr+P$}_=dK zxNa5RmEbC$s={Z*Z`umqzlHAyX;uG&D_r}>d{Cv``She`_+}lRtBx&?Z3?j34aGIP z1+LbK=C!j6y)Unw+B(zJoRc_tx`EB}5P5pyKP6nvdy?ZiN45A*4c7nB!hagLTKruD zwfK)`UTWd~Y=}7r67SD->`Bu{-S1v{GVWf5?vxX^H(E8%KJaRuec}43C(nLhbzjcK z^&dy`8Q0@;AXv@u^mVLR&<_G@Gsa^mPaTuM)^S)->mhKpV~a5k2CK#Y2(WdA9|_kd zpSd3eR`VIe`b$zaDezd0ORO}~tN3fOhJk~y8qfjkwi<~-e>{{Z`Pe`-68 zrsn<>C(r3%=aBl&fLq@x#Aa>IL{m?F{{&l~v8Ui?eQNq8?^)n#-m~H6y{D4*95nUh zJr`_VV@DTx)%5cjBK6y6oOd{e*mxM`7gYh=R>$Y z>dEsFSlyRv;Pu6yope6#|4+bbj;F6RyZ=80Ycs|*k!L=igB>^g3;4mzH~aa2aJA}w z{t|AU>VEzTO+9P-HQ1kJG`~3_P6lr*x$j`s{8qSxMNrM^ABk18T&`D zW1HU`j;*F&#{LOh-OoS6)v}*|1^aS8Yx{+!=6)6@&u?Jokp28S-1@5f`42Sp)b}UY z`i#wfR?{zeInDl(cM7<9tNXbVntJlKfX!=c_OqIP?&s8RFVFqMzWQ@sUY}NS^1PO& zc@OnEwoC;tTfxg!@bVSBLIw9O@PhO!(p<-t>Ae?U$cfRJ8gnmoEdO?ImHYR5tK7fe zTjl=!-YQ?P#{K)fRl9$`x61wdz4EMo7r1raR9x@6!PWeET=!G6UaMG;r4@&uP=bjnPk=Ii0h5Vx|YX&fzn_UF*Eho)PXfWnr4W+T%MDxS5YxS`)!<2bGYz8vEfm1Ddrn))Ps-1Dn}&7q#ptNMVgNt^TWyjM$2tAo`t&Kh9H@j4x! zHT9=ujrxL(Q%~GlVB@?FC$1lwdOkDh4>ndkKQFg7*g5-tMSCm#Nwkc$4w_uQjI}P< zv3zFDSnHvwXRP(X#;W_CNuKvivk=#p=Bp3Q>-}o%Gcx%&EV}*SfT4 zyv@PZp}()_TU7LF9b2MT>(~meR&A+cYxLCNIOf#fI$RI!&d-{+0b5)6wqSiS|Lwr8 zbLPK2T&>zN?hfb~*SzM^-{Y(;$JujYJ({nzX&&eP^y?JZd-J+9kL~*O`B|-<$&q8L z&n^`{5m{QT>nEWT>oQh{PY^XxW;d-@!MqqmYeMn)q_n7uX9||_d4P2kR?+pX1wb6Xel-3`1hIl_m-&6Q`r5n$iT=%+3FXm_x-#O(n#&iQ!UJoa8kJPxDbYWgI83|Kwa zk+EQNs>kPVVAp99dE8^;z-szt4b@y1*T}iqr>5~>Ys!0$31Bs!`?Kz9)pd7`w7J)E zeB5U_K6`;3$KK_4phKcAav3jsmOc zn>AE(U0fsQW}ljl23u2(&oN-N93O3J#%JAKqm2J|u=^u59tT$QK9+mf@o?{9u65#0 z0ITQTcOuwW^~9YFHtr<)#GL|G&v>VTjaAQm_B629Oz*SW-9y#;>_5=t`ng74qtr6S z8DO>4@lUXIoLTWX3#^_x&ITK&p15C%eN7N% zj-zkZOq-v@c;4Oy_T2QP{Z_bI>~|D)$I*T}TrKvy3VWx*ekWWl_InCDm-J@*-Eg(o z?9u8H@ehrsI=wg$Wz>cMS9B zpS67iT-|H`f~%P$>!ud}N5LNJ@PC7i&;ED}tdF|SFY@a1%cj)jYeSmPFB{NrRABGj z8&~iqV6SDH(dSv_N%B>nWuB^V?}0B>xcOeL@B#353hvy@m3xnQa_@N>?6~&cd!D6# zj^@j2hO%2wX8eS>gI&AqyHCJs`et3#JYKGw z^R-VcpMtF=_r%Y@YPlzBQ!_rt!gb5|pMxDg`Mv6?A5X07g1=W3r?z5`oJ=K4KYEpyeTW_-j+H08c6ee_8i;1!0weTDtOBRdtPo;!CM#D_nq6&tbaTDobx|3X7!x^v%*J; z*(6o_Y&G7q!u4Og!mVYM3b*EUE8PDken5q5->t@XukrCUzGsaeRO5%#_^~y9e2w2! z;nsgwg*(58DqQXS;ySk!yKZxLU5|@+oNMGPgLf zQ-X~RpQ_-weoqbe992JsKIi!~aBYd37HpoxO$Rqlz52P>^l)v-GXvQ3KF{hi!hMeN z^Bu>~U!NR%?6^JYvtJi1Z1(Qgh3OZe z`LeD>3%h;Ruou`I*{_R%)v{mZ3(?HwxZ=bv0X8;#$%1F!Ed_VqxmUGi-b;hEC2kq8 zc@noQ*f{s9w!|$5))u}z*n3m<>k9DfSI5v_pB#hUVDBN;?b^wWacx{D=Vza~Rs>sD z_UlSuwd_}IYQ|^nT_@w+uPcKcFZosld;fH7?_Kij*VVx56*1aUb02Vh->iME7&Dyi4>@)6KV8>mBmYn_2)U#jv>rcymT?g#T{i zPM-C_UT4BLfO{PCdBcWqebh6@jlk>Cv`s3Wxi&5|ZO&<3dbRAMf#B>vk6Fgs99>)d zwkZ7csa{vNMAw$Mt-y|v{kAnY`^}vC>ytIy20Vb4{USFe<8BMiKCh0u9lEyoZD085 zQ@wugfUYg|?Fg>cw-Y?|nNxp#Qs2&C>oce8CpRYJ?gF+xzu%~iI|yA{{B|w;^wBr< z3`W-$-U_bPHw2#g%&EUVsc$IQ`mP|S^~sG%zBaJ=vTufg)!Zl94{Grr4%Yu(_K)xN zMu63_e@B9Sxqr3oK^sGJ|B5r;QDE1_>u&Y9jYikz*t^rqJ-^Ir4s#{PSa9~Hc~axw z(6uFpyqd#Y@f`vRuS_rSjB z+8ldNdbxF)*Bs_bj{U*rFur;o9)PYbIpozG=E{5y1Urtto8;aBz-EbsryruFbIzp_f~ydCg(2J&0Z&+uy;}xgG~s%Un+c`*N<@PN1Davrcj9JPDj*;`x>QC!=d~?BnU>)@fdI zm@7F>1)Ib8>fSgFU0ZUd(AEU?#^&NOYgexF_FOPe2T9({BDIR`vI4ZHQpjY)myg00WvTs_w3p=(P$=Y!p| z;TM3b^<4;0ecDprMTKsC<}{DKsqbR2^*Ofn$&E>Umwe-pi0;%@`1 z<$iNJSnVWo_g`+{~Ij3U^JlWgY$vR*%nP;D>1D zcaCy>oP+bR4*QJzIM_Ne?h{}&$ITk48J{(BJ{kW>u*V_K2~UC5JPtW-YQ}p^9N*)h z?HT&4|Feb7-t~W;{so#ZV_z)n_No0PusO2+FN4*x{_^K&=5kzdVqXCp8~$p+vyT6P zyVmNN>uX^3_`D8wtUDcCMlM`m-)? zfgL~fy$$wQm#u7 z>Z#{raP@uCCvY`=9bcQ8^IM1B^|H^npMo8?`tzZmp{eJ6(&u2c+@t>o_T_o6?F*Wk zxy7mND{%Gt^EJGB{rLv2k9y|#EqFbeHtYYAUM<($@4>kzRPU)jplgfYkA+|TIp!yH zZK?5Ruw&$!`wQ4>vDaL4>aS1M>sRo1wCpFjF&Xzau=ROQvt^v$(Y3|zkHXLUnk{~R zqH7EH1JG)HQ@~T7IrY~k^>qSUpE+GWxiJ~H1#ErRUmdqIy0-XrDg2y6_1@MMU0dqw z2CmjOB|P<+Q-6I@-&A1h%X=cZG08VI*nGLxO#@eRPh>x+#ecfOKliBV;cD5dGl6}% zSGCPZo0aBX6=%LPgIyQvsvfsl(6u?XpFhhzm&|Jpb0x=Y;Ox!nd!pIVwIzqVn!{X~ z&m3UK(bvzv<+05Lc5O1(x#4P=>wI8e&Q;sIwC*(P6sOMl!Pe<COoU69pv{h)vz9GUz+^pG>^Wy{`3cX{c(QQCpRYbtqrz5=UzS5>!52(J?nzqv*GK3tM#o9 zPkq`_-v)(ledaWezNv3Ru=P2%^~sG%eH(#ujH}0IV{~n)XA`jdEPPXNwZ6^ZsZU$# z8&K%hXHN6zoB9TVtDdn0@+aJ9az;i*qs>f5H! zt=?{oT4`Zk&`&$(fR-QF5{qVEPa zN1oG%gVpk!E+0ZOm*a{PyF1v}@I4BiXXue|HIHf5c@*4r)|Pb`4OWlO81M+1`JJO& zALrnFtiwLzjs;ss#{C;u&2h6vYQ|@coKMCd2lhDRSdIs)c^q=w)QtC-IKIb0+eG@T z|DJ`--u2&`ejl1IWA`oW_Nje8usO2+`-9c8{_?$O=5kzdVh;ct8-8HHvyKPBU2FBs z^5= zhl4#9d9FV~Osl?6IufpKy+_f@^FHZl@OnjFZJFmW;QITOW6{(z*5AR#tEZmhz}5Fj z$HUe1b$o4V&M)s}>=S>2Hhlb96mfnQr4uy{qWRj1=Fg(+K)-W=cckBi=KAkS@6VR_ zvjQiRgMTGIqsGsz@hfWl<{H1R#-FJ1=W6`z8h^jWKdSN1D%?4JS>wOec&8~R&u=RH zyVdwqH9lK~yFNW?e6b4m=XI8=aQ{BS#uW~1Y*yjcyKRl{TjNL8_;EFUa*dx+;nsUz zjbB#bJ+NO__+Nz^|3ig){C}-*?JbCGneu0?rmk@9Gt~GjH9k*;0~-rgxbaI=xc22MT>Hv3 z-lxX<)%b=L?%$i(qQc#etrZSz46X5D74G=E*Z7zk->bs?*|vi!9N0La#?P*B^PgMe z=U2G-E~{{0028O^shy;m-et3it1C+*accRk;3-RXDKmRE4uu8ZTA2_E#$0 z_^%7@eZlklR5p{>-26P%eBjg2)cqX|@_&F&p{eiBb=c>o)4}TA1H2~7lk=Zob55dp zA27#RXzIy%Huy}MdUBovRyU{Db-D45cOKYr#=yC~G|mUBO)8!>F956gd53k3!{;Kn zKOcV~eeOpWgVl4-y98{kdY*qS1*@AY@0%`zlZyWWj_*9c`ONZiur_P<-Xl-VSAv~K zYQ74rmYS~yt5s{h7H-Yg(5L3>!0M^_da$wTsrd%5dbQ>o;i*}BYQ71q&6>SG%2V?# zVAna%H@AY-x^fTCJZ^)lXCAkM)vEKj172VIJJHlLkGsIes%IW|gVn3^xCib$+|SxG zk9)z|oQL;ydFF9H*tO3*9t5jBK+8NH0;^{p4};aJ^LPYaKOX-=Q_nme1skiLdHfr! zUY*Bd@XSMd=J7aKoAdB_M4ov(33hJbPl4TMhp@*zFP;YLqn^1w16KDLI)2Z>)$^R{ zbB&TNq?Wv|fE_!(|K?S&n*Qlzua7XR15j+2~kfYquw-zHdfu*eXpXH+GhuQO-b!@z|~UwoM6{7wa*3DM?JOA z4OUML^MI?h&kNT_J+;pdws!wEUE&r1tEcwvU}M#--S<^$V&7Ap#TvMWdxG6l_U`e8 z=ohB>(!NMxx6iy61)C$Ear6SKdCt0D}z zpK+G~no{{6j`X#VZJtk=q5k7shP0ycN-(h- zntIlJb+B>j*?Vh%9ZOrrUK8x`$k=_su4l$x3#^v0`+@7nyT5i?#$Fq2oO;Gy2kcnd z{QXkv(yvFeCT-g_+W(Ht`e4^bo1bH}(mYKxc?d)8;<%cf<_^?@BY?ymEb#w+(FG z)G-XKHi=vrYd3htat$)paCB`MYXsP_%;8vaeKOYW;Oba=z|~xrj5QMMSdLSj*C;e? z8EZ7yvCQFEa(yz^7;ts0v2ZoV%3S{jR1XbJ z!D{9{fZp2dy#^mde=yCL_DO}^KKTy;n?vsN`k`Q-v2qVQ46c^vo5R6sXW*OXntD91pIJaRR(L z#))vXJl~uIHb&hwJ&s;2`Az{lM#eu4tad6bKK}q4mpyPgSU+`ho=mS6|1-hnjQ>Bu zYVOyJa~53PeVw?o!PRljf$OK9T<3z-&E?oAI=AReC*vX7t&uu^QHab z!fx;Q?vG2r=Eyz!Qm~re33}}17tqXQZn5KDMxSvnFKqUXdnNrt~jyRgRMFI2C%v0Uej*`yGF6!1lHf<;GAy(tKCe?KD`yJo;AM> zY##OGxgG2{iMa!;mY6%i#;NDpa~Ifak2dpo-M<^Gp6lp6;Ol74%elC|YWlf;uBB^W zpIYt(TZ>~mr~ANap2s;pYR2dIxR!GFg!$}~?|!grlzb0>)sjz}n(^*0-q92T+P_e=;c1^cEa{K*xxa&{{QBl z{w03q&i&^rxY{YjebDi~hO4KZZ@|XnzW6O%ts3(kTs^;g_It2-)UEq}^lIjGj32=# zGp1|ocL6`a^*In6pP%9CInRFq8>ep0AL!NM{~K5>?|ps;t4(4IuVsIL{mfC_aeeRp zC)oEx7gCFP^z*}1^}HXN0&c8&KEvz;RyUuu{EA&IIXZ*Y@)>y-{YhaOZK=5{*v~9d zb2qqJ#+edqta?5ZoeHe(xPIp)&$!co9XEX1f~St@;Hg7f#+e>$takJI-I-eQ%m`M? zn#=@Ndx!a_u9@NL#`|5HTofmvYWuMIlcc0~Y;a;2{O+EW;0kCoE$ z+*obS+wT+Al4BvTT8{O?V72VeMZoSS_3Y0@!S2s%D*Lk+ntG1KVqjy{b37LZt7m`u zU8b6~xW|_St7Ts-1`vMN}ux<6Ngs~f*Ey*%?=9c-;RzH5Tj*08Z= z?vuV?&$sN8wZQtRo73-mn;N(uID5W7*nMg5o?nN4U79cV{Cb7mK4Yy9Hb?gS24J;( zJ|$n9W-fD!tzkp@jJr``vv=H0=r^Th+|3HRea0OCHb=%C2v*Cu@{MWca$Ip@w**^r z_*P(Zg>MaZjl#D9>z_5>7Odu)XN|W58>cOC+k?#;pB=#Zse9h|47MZKd$u;?{7h*l zu)5E(e&4$@cng|!TdV7*rk}^fwQ`R38G9FSb?iZKwLDj7Q!_s6?^?+{=H|0czFont zLFO_Ttd@M*)QrzQ@|YX%??oK~b`EmS>o)MzH0LsuULME5VIvT8=_kqWN529JOwYq+4`gvSjE9YpR zv5y5;$NoE9Eqh&?n(Ay2s)KdNuLc^sPng2K48^J$6~g^T0FE^gWke9@_=rl?&VX^m4Bm z&iO*HzXQ*F#%a&IE&^La?#CB{^;eJ2C17it#PKlSrC>FEQ@5J_)_pnHy0u+KFHf#3 z!BZEuE9m91T@7Bouw6wj*XDKUT5#rm4OkxA_2A6?IR`MUH^z^x(QSv?8XN4*pE zc&?b|F}OC@mag$M|EE^HUOWR=kI%D(PxadH9GZH3o-cf=*MS$%)CaPIc-(Hh2sXE| z+N|d#di8wI^)h%KeATr-MX%;OJ)d3$yT5X5{sUI?dKRD8;Kt=T@;X>Qb#uN#uNMC| zz-pPxo8bDIyoIJdsW`9S2CEsHxOc#ba~!#T8RuPab)5GK|LVHFkEWh+J^-s3n{oaN zb{yjzN3LJS`4C(k=OeiO)noH9Ts`A_0#-9N<9rHEoa4y#bDYoU|3~|x;O@aM!R}Mv zGx@iwz5=W1@Ac&yu=&jOHND*9uI*d!9Q5i#@He;d>fY<-p!Z&BZ$8hl@4&Ora(hdrBuffer.Offset, hdrParams); _pipeline.SetUniformBuffers([new BufferAssignment(1, buffer.Range), new BufferAssignment(2, hdrBuffer.Range)]); diff --git a/src/Ryujinx.Graphics.Vulkan/Ryujinx.Graphics.Vulkan.csproj b/src/Ryujinx.Graphics.Vulkan/Ryujinx.Graphics.Vulkan.csproj index e59823648..19922e13b 100644 --- a/src/Ryujinx.Graphics.Vulkan/Ryujinx.Graphics.Vulkan.csproj +++ b/src/Ryujinx.Graphics.Vulkan/Ryujinx.Graphics.Vulkan.csproj @@ -16,6 +16,8 @@ + + diff --git a/src/Ryujinx.Graphics.Vulkan/Shaders/ColorBlitHdrFragmentShaderSource.frag b/src/Ryujinx.Graphics.Vulkan/Shaders/ColorBlitHdrFragmentShaderSource.frag index 6a8db28fc..35b9010e7 100644 --- a/src/Ryujinx.Graphics.Vulkan/Shaders/ColorBlitHdrFragmentShaderSource.frag +++ b/src/Ryujinx.Graphics.Vulkan/Shaders/ColorBlitHdrFragmentShaderSource.frag @@ -8,6 +8,8 @@ layout (std140, binding = 2) uniform hdr_params vec4 hdr_params_data; // highlight colour mix: 0 = luma (calm/white highlights), 1 = max channel (coloured lights pop) float hdr_blend_data; + // highlight whitening: 0 = keep colour (old look), 1 = fully desaturate extreme highlights to white + float hdr_whiten_data; }; layout (location = 0) in vec2 tex_coord; @@ -41,8 +43,15 @@ vec3 sdrToHdr(vec3 srgb) float luma = dot(lin, vec3(0.2126, 0.7152, 0.0722)); float maxc = max(max(lin.r, lin.g), lin.b); float hl = mix(luma, maxc, hdr_blend_data); - float boost = mix(1.0, peak / paperWhite, pow(clamp(hl, 0.0, 1.0), curve)); - return lin * paperWhite * boost; + float hlAmount = pow(clamp(hl, 0.0, 1.0), curve); + float boost = mix(1.0, peak / paperWhite, hlAmount); + vec3 outc = lin * paperWhite * boost; + // Desaturate the most extreme highlights toward white: a blinding light (sun, fire, a bright lamp) + // physically reads as white, not a saturated colour. Luminance-preserving (mix toward a grey of + // equal brightness) so only the harsh hue at the very top is removed, not the brightness itself. + float t = hdr_whiten_data * hlAmount; + float lumaOut = dot(outc, vec3(0.2126, 0.7152, 0.0722)); + return mix(outc, vec3(lumaOut), t); } void main() diff --git a/src/Ryujinx.Graphics.Vulkan/Shaders/SpirvBinaries/ColorBlitHdrFragment.spv b/src/Ryujinx.Graphics.Vulkan/Shaders/SpirvBinaries/ColorBlitHdrFragment.spv index b86284304ba8122bf60b3b71574e37304661a472..ff6bc4fdca878c395902bab61fb90e27074c69e7 100644 GIT binary patch literal 3520 zcmZ9N*>+S_5QaA-0b~%6NffaIs33?SLx9K-6@wZK2pUnb>2x=wrIU`GZU#q!qIlzt zUicoqfluHI_#WQq^7~Go3dfwaYVG=~{u*}eea`7yx^sC-OVX;eCjFFp&-%0!OTxA+ zTSvz(jqRW7)b<}fazMu8sV^hSS(}!pepY3=(wNcl6*7;kgQuMWXnkZ&=?{YX)5VW@Q!m#V?RvG-xK)?8qQGm+RGZyekuU(GSa*O*Yp5|- zYhP}Sk*w0*ee3$6;p+#7(*X5X;j1FXt!BMZnVzm>cv)J{n#G-~^@4KVI_>QnlPp=> zHr7XR18F;Ws@5*gR@#;6xx~blXEynd?B!aeQ(5piH+_@M`b@3oTJSj|eYd9?o%&4A zY40WMKI*$!nN`-$9`vks=Ic*;|Azm>0+jjOt-o-ou0oxO`uG+T5}ql^;OnZx3dth621ny zwZ^-hEVuUF8TF+*^6Gc8I^HSta<$cJ*U-&7FSpffb=y6B9~`q7rBBbsd7guiPEO~?U+n2{|Mq7qn zLbky7{SoUgfmg6MNt<&Nd$)W%i!WHsXMXu#vYJb~e6e>;oBn&O z&d~GHFP=xV9px;|(~tOQ`y0yut6c9+2l0*tzMtch`0ep7Vt-m=TYQ|{`o2f}W^2EQ&AB{4wEKkZ2Xt-L3EPjk z&AJ!JX0HzsF-7w7pnYwTwOx}4Qy zFvn)ZzIG$}eE;OSmo11nE`a1dncKAYe7zi7!RB}hF`vADsjI!%ue0*;I}rW)&9lcR zSUqF=-iBB&&fMOLcuxV#_kR<9vHzPQ_Fq%PM;GA}i}1)I?7yZM@BgNVCkogzxLUyW zQ(c4`IUZrIepjAHO2`gm7ZUS-0bQSWz?q2S9A89tcKY_PiX(0xx_o`kMtnR!!)g#SwEDy@VKZ2#J^@==zND9K;cG484RHGlWFU zaddsgcvfP4y_xs!>aY3xb_U-t;3+AFt>3;JlG@ihZBOrrr0% zci~NB31Tke^+)_@ZZGcP95(F{{}y`0n@hj_jZhQFXBE#_xuB2LASq$`8j@sE=PNCK1MfB^!5pQ?8}~vkz+mILH)*hPd-IFtH@PG zx8Bdx(|-+-qdhnkbn^sf68##Y-5S2v+Ttwg=q03vM2_p|+5_J}cMlKpeNLgv(QeKv zt2X)e^%-KFxZgK(TiovkSX=O$=<=i9G`gJlX3wB&3khBeo3(>Ci>@uincL6jNX*cfxaIqFyRZPpd&WUVnp|xb&-rtoy+4&cuL6mm}$=i4tbY!!>fFha^>M-ln+X^H2HUB+(xXOG~($VYmrt_ z!%mttgIn@3OE_bS_hp|5ETsOMY#_i>PmBmS@;_xmp)8Om=JVj@^>Lb$+{fi!OR=sr zttrSJ`S`7R(6T3>gP2_h0(&+zgV~1yHSlYC6=wYq-~k6s=w9~EL7sW4?b$B{h#<*w z;zbQ-kew1pG?==mvrvm-dldtJVN*8l8r^8qIBpCiL zoaK5S6we8$Wa&p?M(BGHpX5^mpVh3(egxja z6Wj{T6lBA;d=e21K>!Wh8h<=_f2HA>P=OL diff --git a/src/Ryujinx.Graphics.Vulkan/Window.cs b/src/Ryujinx.Graphics.Vulkan/Window.cs index f947f0a0e..21208a2ab 100644 --- a/src/Ryujinx.Graphics.Vulkan/Window.cs +++ b/src/Ryujinx.Graphics.Vulkan/Window.cs @@ -48,6 +48,7 @@ namespace Ryujinx.Graphics.Vulkan private float _hdrCurve = 2.8f; private float _hdrGamma = 2.2f; private float _hdrBlend = 0.5f; + private float _hdrWhiten = 0.4f; public unsafe Window(VulkanRenderer gd, SurfaceKHR surface, PhysicalDevice physicalDevice, Device device) { @@ -450,7 +451,7 @@ namespace Ryujinx.Graphics.Vulkan int dstY0 = crop.FlipY ? dstPaddingY : _height - dstPaddingY; int dstY1 = crop.FlipY ? _height - dstPaddingY : dstPaddingY; - if (_scalingFilter != null) + if (_scalingFilter != null && _scalingFilter.IsResolutionSupported(view.Width, view.Height, _width, _height)) { _scalingFilter.Run( view, @@ -466,7 +467,8 @@ namespace Ryujinx.Graphics.Vulkan _hdrPeak / 80f, _hdrCurve, _hdrGamma, - _hdrBlend + _hdrBlend, + _hdrWhiten ); } else @@ -485,7 +487,8 @@ namespace Ryujinx.Graphics.Vulkan _hdrPeak / 80f, _hdrCurve, _hdrGamma, - _hdrBlend); + _hdrBlend, + _hdrWhiten); } Transition( @@ -620,6 +623,15 @@ namespace Ryujinx.Graphics.Vulkan _scalingFilter = new AreaScalingFilter(_gd, _device); } + break; + case ScalingFilter.Nis: + if (_scalingFilter is not NisScalingFilter) + { + _scalingFilter?.Dispose(); + _scalingFilter = new NisScalingFilter(_gd, _device); + } + + _scalingFilter.Level = _scalingFilterLevel; break; } } @@ -687,13 +699,14 @@ namespace Ryujinx.Graphics.Vulkan _swapchainIsDirty = true; } - public override void SetHdrMode(bool enabled, float paperWhite, float peak, float curve, float gamma, float blend) + public override void SetHdrMode(bool enabled, float paperWhite, float peak, float curve, float gamma, float blend, float whiten) { _hdrPaperWhite = paperWhite; _hdrPeak = peak; _hdrCurve = curve; _hdrGamma = gamma; _hdrBlend = blend; + _hdrWhiten = whiten; if (_hdrEnabled != enabled) { diff --git a/src/Ryujinx.Graphics.Vulkan/WindowBase.cs b/src/Ryujinx.Graphics.Vulkan/WindowBase.cs index 99b3864e5..d36935633 100644 --- a/src/Ryujinx.Graphics.Vulkan/WindowBase.cs +++ b/src/Ryujinx.Graphics.Vulkan/WindowBase.cs @@ -16,6 +16,6 @@ namespace Ryujinx.Graphics.Vulkan public abstract void SetScalingFilter(ScalingFilter scalerType); public abstract void SetScalingFilterLevel(float scale); public abstract void SetColorSpacePassthrough(bool colorSpacePassthroughEnabled); - public abstract void SetHdrMode(bool enabled, float paperWhite, float peak, float curve, float gamma, float blend); + public abstract void SetHdrMode(bool enabled, float paperWhite, float peak, float curve, float gamma, float blend, float whiten); } } diff --git a/src/Ryujinx/Systems/AppHost.cs b/src/Ryujinx/Systems/AppHost.cs index 80e734718..7ef92e36b 100644 --- a/src/Ryujinx/Systems/AppHost.cs +++ b/src/Ryujinx/Systems/AppHost.cs @@ -212,6 +212,7 @@ namespace Ryujinx.Ava.Systems ConfigurationState.Instance.Graphics.HdrIntensity.Event += UpdateHdrLevel; ConfigurationState.Instance.Graphics.HdrGamma.Event += UpdateHdrLevel; ConfigurationState.Instance.Graphics.HdrHighlightMix.Event += UpdateHdrLevel; + ConfigurationState.Instance.Graphics.HdrHighlightWhiten.Event += UpdateHdrLevel; ConfigurationState.Instance.Graphics.VSyncMode.Event += UpdateVSyncMode; ConfigurationState.Instance.Graphics.CustomVSyncInterval.Event += UpdateCustomVSyncIntervalValue; ConfigurationState.Instance.Graphics.EnableCustomVSyncInterval.Event += UpdateCustomVSyncIntervalEnabled; @@ -323,7 +324,8 @@ namespace Ryujinx.Ava.Systems ConfigurationState.Instance.Graphics.HdrPeakBrightness, HdrIntensityToCurve(ConfigurationState.Instance.Graphics.HdrIntensity), ConfigurationState.Instance.Graphics.HdrGamma / 100f, - ConfigurationState.Instance.Graphics.HdrHighlightMix / 100f); + ConfigurationState.Instance.Graphics.HdrHighlightMix / 100f, + ConfigurationState.Instance.Graphics.HdrHighlightWhiten / 100f); } // Maps the user-facing HDR intensity (0-100, higher = more pop) to the highlight curve @@ -695,6 +697,7 @@ namespace Ryujinx.Ava.Systems ConfigurationState.Instance.Graphics.HdrIntensity.Event -= UpdateHdrLevel; ConfigurationState.Instance.Graphics.HdrGamma.Event -= UpdateHdrLevel; ConfigurationState.Instance.Graphics.HdrHighlightMix.Event -= UpdateHdrLevel; + ConfigurationState.Instance.Graphics.HdrHighlightWhiten.Event -= UpdateHdrLevel; _topLevel.PointerMoved -= TopLevel_PointerEnteredOrMoved; _topLevel.PointerEntered -= TopLevel_PointerEnteredOrMoved; @@ -1156,7 +1159,8 @@ namespace Ryujinx.Ava.Systems ConfigurationState.Instance.Graphics.HdrPeakBrightness, HdrIntensityToCurve(ConfigurationState.Instance.Graphics.HdrIntensity), ConfigurationState.Instance.Graphics.HdrGamma / 100f, - ConfigurationState.Instance.Graphics.HdrHighlightMix / 100f); + ConfigurationState.Instance.Graphics.HdrHighlightMix / 100f, + ConfigurationState.Instance.Graphics.HdrHighlightWhiten / 100f); Width = (int)RendererHost.Bounds.Width; Height = (int)RendererHost.Bounds.Height; diff --git a/src/Ryujinx/Systems/Configuration/ConfigurationFileFormat.cs b/src/Ryujinx/Systems/Configuration/ConfigurationFileFormat.cs index 92e8da70a..4a7e5a1b1 100644 --- a/src/Ryujinx/Systems/Configuration/ConfigurationFileFormat.cs +++ b/src/Ryujinx/Systems/Configuration/ConfigurationFileFormat.cs @@ -17,7 +17,7 @@ namespace Ryujinx.Ava.Systems.Configuration /// /// The current version of the file format /// - public const int CurrentVersion = 77; + public const int CurrentVersion = 78; /// /// Version of the configuration file format @@ -99,6 +99,11 @@ namespace Ryujinx.Ava.Systems.Configuration /// public int HdrHighlightMix { get; set; } + /// + /// HDR highlight whitening (0-100): desaturate the brightest highlights toward white. 0 = old look (keep colour). + /// + public int HdrHighlightWhiten { get; set; } + /// /// Dumps shaders in this local directory /// diff --git a/src/Ryujinx/Systems/Configuration/ConfigurationState.Migration.cs b/src/Ryujinx/Systems/Configuration/ConfigurationState.Migration.cs index cf89374b1..df3ddb0c7 100644 --- a/src/Ryujinx/Systems/Configuration/ConfigurationState.Migration.cs +++ b/src/Ryujinx/Systems/Configuration/ConfigurationState.Migration.cs @@ -91,6 +91,7 @@ namespace Ryujinx.Ava.Systems.Configuration Graphics.HdrIntensity.Value = cff.HdrIntensity; Graphics.HdrGamma.Value = cff.HdrGamma; Graphics.HdrHighlightMix.Value = cff.HdrHighlightMix; + Graphics.HdrHighlightWhiten.Value = cff.HdrHighlightWhiten; Graphics.VSyncMode.Value = cff.VSyncMode; Graphics.EnableCustomVSyncInterval.Value = cff.EnableCustomVSyncInterval; Graphics.CustomVSyncInterval.Value = cff.CustomVSyncInterval; @@ -554,7 +555,8 @@ namespace Ryujinx.Ava.Systems.Configuration }), (75, static cff => cff.HdrIntensity = 80), (76, static cff => cff.HdrGamma = 220), - (77, static cff => cff.HdrHighlightMix = 50) + (77, static cff => cff.HdrHighlightMix = 50), + (78, static cff => cff.HdrHighlightWhiten = 0) ); } } diff --git a/src/Ryujinx/Systems/Configuration/ConfigurationState.Model.cs b/src/Ryujinx/Systems/Configuration/ConfigurationState.Model.cs index e75f30265..5af922bbb 100644 --- a/src/Ryujinx/Systems/Configuration/ConfigurationState.Model.cs +++ b/src/Ryujinx/Systems/Configuration/ConfigurationState.Model.cs @@ -652,6 +652,12 @@ namespace Ryujinx.Ava.Systems.Configuration /// public ReactiveObject HdrHighlightMix { get; private set; } + /// + /// HDR highlight whitening (0-100): desaturate the brightest highlights toward white. + /// 0 = keep full colour (old look), 100 = strong whitening of extreme highlights. + /// + public ReactiveObject HdrHighlightWhiten { get; private set; } + /// /// Preferred GPU /// @@ -706,6 +712,8 @@ namespace Ryujinx.Ava.Systems.Configuration HdrGamma.LogChangesToValue(nameof(HdrGamma)); HdrHighlightMix = new ReactiveObject(); HdrHighlightMix.LogChangesToValue(nameof(HdrHighlightMix)); + HdrHighlightWhiten = new ReactiveObject(); + HdrHighlightWhiten.LogChangesToValue(nameof(HdrHighlightWhiten)); } } diff --git a/src/Ryujinx/Systems/Configuration/ConfigurationState.cs b/src/Ryujinx/Systems/Configuration/ConfigurationState.cs index 59073427b..5a947a062 100644 --- a/src/Ryujinx/Systems/Configuration/ConfigurationState.cs +++ b/src/Ryujinx/Systems/Configuration/ConfigurationState.cs @@ -46,6 +46,7 @@ namespace Ryujinx.Ava.Systems.Configuration HdrIntensity = Graphics.HdrIntensity, HdrGamma = Graphics.HdrGamma, HdrHighlightMix = Graphics.HdrHighlightMix, + HdrHighlightWhiten = Graphics.HdrHighlightWhiten, GraphicsShadersDumpPath = Graphics.ShadersDumpPath, LoggingEnableDebug = Logger.EnableDebug, LoggingEnableStub = Logger.EnableStub, @@ -220,6 +221,7 @@ namespace Ryujinx.Ava.Systems.Configuration Graphics.HdrIntensity.Value = 80; Graphics.HdrGamma.Value = 220; Graphics.HdrHighlightMix.Value = 50; + Graphics.HdrHighlightWhiten.Value = 0; System.EnablePtc.Value = true; System.EnableInternetAccess.Value = false; System.EnableFsIntegrityChecks.Value = true; diff --git a/src/Ryujinx/UI/ViewModels/SettingsViewModel.cs b/src/Ryujinx/UI/ViewModels/SettingsViewModel.cs index 58be9ea2a..da122ee83 100644 --- a/src/Ryujinx/UI/ViewModels/SettingsViewModel.cs +++ b/src/Ryujinx/UI/ViewModels/SettingsViewModel.cs @@ -446,6 +446,19 @@ namespace Ryujinx.Ava.UI.ViewModels public string HdrHighlightMixText => $"{HdrHighlightMix}"; + public string HdrHighlightWhitenText => $"{HdrHighlightWhiten}"; + + public int HdrHighlightWhiten + { + get; + set + { + field = value; + OnPropertyChanged(); + OnPropertyChanged(nameof(HdrHighlightWhitenText)); + } + } + public int HdrHighlightMix { get; @@ -850,6 +863,7 @@ namespace Ryujinx.Ava.UI.ViewModels HdrIntensity = config.Graphics.HdrIntensity.Value; HdrGamma = config.Graphics.HdrGamma.Value; HdrHighlightMix = config.Graphics.HdrHighlightMix.Value; + HdrHighlightWhiten = config.Graphics.HdrHighlightWhiten.Value; EnableMacroHLE = config.Graphics.EnableMacroHLE; EnableColorSpacePassthrough = config.Graphics.EnableColorSpacePassthrough; ResolutionScale = config.Graphics.ResScale == -1 ? 4 : config.Graphics.ResScale - 1; @@ -986,6 +1000,7 @@ namespace Ryujinx.Ava.UI.ViewModels config.Graphics.HdrIntensity.Value = HdrIntensity; config.Graphics.HdrGamma.Value = HdrGamma; config.Graphics.HdrHighlightMix.Value = HdrHighlightMix; + config.Graphics.HdrHighlightWhiten.Value = HdrHighlightWhiten; config.Graphics.EnableMacroHLE.Value = EnableMacroHLE; config.Graphics.EnableColorSpacePassthrough.Value = EnableColorSpacePassthrough; config.Graphics.ResScale.Value = ResolutionScale == 4 ? -1 : ResolutionScale + 1; diff --git a/src/Ryujinx/UI/Views/Settings/SettingsGraphicsView.axaml b/src/Ryujinx/UI/Views/Settings/SettingsGraphicsView.axaml index d05741a14..fa60ea49f 100644 --- a/src/Ryujinx/UI/Views/Settings/SettingsGraphicsView.axaml +++ b/src/Ryujinx/UI/Views/Settings/SettingsGraphicsView.axaml @@ -209,6 +209,28 @@ VerticalAlignment="Center" Text="{Binding HdrHighlightMixText}"/> + + + + + +