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 8ca5e52bb..3b106f7d7 100644 Binary files a/src/Ryujinx.Graphics.Vulkan/Effects/Shaders/FsrSharpeningHdr.spv and b/src/Ryujinx.Graphics.Vulkan/Effects/Shaders/FsrSharpeningHdr.spv differ diff --git a/src/Ryujinx.Graphics.Vulkan/Effects/Shaders/NisScaler.h b/src/Ryujinx.Graphics.Vulkan/Effects/Shaders/NisScaler.h new file mode 100644 index 000000000..1ca7d4a04 --- /dev/null +++ b/src/Ryujinx.Graphics.Vulkan/Effects/Shaders/NisScaler.h @@ -0,0 +1,1023 @@ +// 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 SDK - v1.0.3 +//--------------------------------------------------------------------------------- +// The NVIDIA Image Scaling SDK provides a single spatial scaling and sharpening algorithm +// for cross-platform support. The scaling algorithm uses a 6-tap scaling filter combined +// with 4 directional scaling and adaptive sharpening filters, which creates nice smooth images +// and sharp edges. In addition, the SDK provides a state-of-the-art adaptive directional sharpening algorithm +// for use in applications where no scaling is required. +// +// The directional scaling and sharpening algorithm is named NVScaler while the adaptive-directional-sharpening-only +// algorithm is named NVSharpen. Both algorithms are provided as compute shaders and +// developers are free to integrate them in their applications. Note that if you integrate NVScaler, you +// should NOT integrate NVSharpen, as NVScaler already includes a sharpening pass +// +// Pipeline Placement +// ------------------ +// The call into the NVIDIA Image Scaling shaders must occur during the post-processing phase after tone-mapping. +// Applying the scaling in linear HDR in-game color-space may result in a sharpening effect that is +// either not visible or too strong. Since sharpening algorithms can enhance noisy or grainy regions, it is recommended +// that certain effects such as film grain should occur after NVScaler or NVSharpen. Low-pass filters such as motion blur or +// light bloom are recommended to be applied before NVScaler or NVSharpen to avoid sharpening attenuation. +// +// Color Space and Ranges +// ---------------------- +// NVIDIA Image Scaling shaders can process color textures stored as either LDR or HDR with the following +// restrictions: +// 1) LDR +// - The range of color values must be in the [0, 1] range +// - The input color texture must be in display-referred color-space after tone mapping and OETF (gamma-correction) +// has been applied +// 2) HDR PQ +// - The range of color values must be in the [0, 1] range +// - The input color texture must be in display-referred color-space after tone mapping with Rec.2020 PQ OETF applied +// 3) HDR Linear +// - The recommended range of color values is [0, 12.5], where luminance value (as per BT. 709) of +// 1.0 maps to brightness value of 80nits (sRGB peak) and 12.5 maps to 1000nits +// - The input color texture may have luminance values that are either linear and scene-referred or +// linear and display-referred (after tone mapping) +// +// If the input color texture sent to NVScaler/NVSharpen is in HDR format set NIS_HDR_MODE define to either +// NIS_HDR_MODE_LINEAR (1) or NIS_HDR_MODE_PQ (2). +// +// Supported Texture Formats +// ------------------------- +// Input and output formats: +// Input and output formats are expected to be in the rages defined in previous section and should be +// specified using non-integer data types such as DXGI_FORMAT_R8G8B8A8_UNORM. +// +// Coefficients formats: +// The scaler coefficients and USM coefficients format should be specified using float4 type such as +// DXGI_FORMAT_R32G32B32A32_FLOAT or DXGI_FORMAT_R16G16B16A16_FLOAT. +// +// Resource States, Buffers, and Sampler: +// The game or application calling NVIDIA Image Scaling SDK shaders must ensure that the textures are in +// the correct state. +// - Input color textures must be in pixel shader read state. Shader Resource View (SRV) in DirectX +// - The output texture must be in read/write state. Unordered Access View (UAV) in DirectX +// - The coefficients texture for NVScaler must be in read state. Shader Resource View (SRV) in DirectX +// - The configuration variables must be passed as constant buffer. Constant Buffer View (CBV) in DirectX +// - The sampler for texture pixel sampling. Linear clamp SamplerState in Direct +// +// Adding NVIDIA Image Scaling SDK to a Project +// -------------------------------------------- +// Include NIS_Scaler.h directly in your application or alternative use the provided NIS_Main.hlsl shader file. +// Use NIS_Config.h to get the ideal shader dispatch values for your platform, to configure the algorithm constant +// values (NVScalerUpdateConfig, and NVSharpenUpdateConfig), and to access the algorithm coefficients (coef_scale and coef_USM). +// +// Defines: +// NIS_SCALER: default (1) NVScaler, (0) fast NVSharpen only, no upscaling +// NIS_HDR_MODE: default (0) disabled, (1) Linear, (2) PQ +// NIS_BLOCK_WIDTH: pixels per block width. Use GetOptimalBlockWidth query for your platform +// NIS_BLOCK_HEIGHT: pixels per block height. Use GetOptimalBlockHeight query for your platform +// NIS_THREAD_GROUP_SIZE: number of threads per group. Use GetOptimalThreadGroupSize query for your platform +// NIS_USE_HALF_PRECISION: default (0) disabled, (1) enable half pression computation +// NIS_HLSL: (1) enabled, (0) disabled +// NIS_HLSL_6_2: default (0) HLSL v5, (1) HLSL v6.2 forces NIS_HLSL=1 +// NIS_GLSL: (1) enabled, (0) disabled +// NIS_VIEWPORT_SUPPORT: default(0) disabled, (1) enable input/output viewport support +// NIS_NV12_SUPPORT: default(0) disabled, (1) enable NV12 input +// NIS_CLAMP_OUTPUT: default(0) disabled, (1) enable output clamp +// +// Default NVScaler shader constants: +// [NIS_BLOCK_WIDTH, NIS_BLOCK_HEIGHT, NIS_THREAD_GROUP_SIZE] = [32, 24, 256] +// +// Default NVSharpen shader constants: +// [NIS_BLOCK_WIDTH, NIS_BLOCK_HEIGHT, NIS_THREAD_GROUP_SIZE] = [32, 32, 256] +// +// NIS_UNROLL: default [unroll] +// NIS_UNROLL_INNER: default NIS_UNROLL, define in case of a compiler error for inner nested loops +//--------------------------------------------------------------------------------- + +// NVScaler enable by default. Set to 0 for NVSharpen only +#ifndef NIS_SCALER +#define NIS_SCALER 1 +#endif + +// HDR Modes +#define NIS_HDR_MODE_NONE 0 +#define NIS_HDR_MODE_LINEAR 1 +#define NIS_HDR_MODE_PQ 2 +#ifndef NIS_HDR_MODE +#define NIS_HDR_MODE NIS_HDR_MODE_NONE +#endif +#define kHDRCompressionFactor 0.282842712f + +// Viewport support +#ifndef NIS_VIEWPORT_SUPPORT +#define NIS_VIEWPORT_SUPPORT 0 +#endif + +// HLSL, GLSL +#if NIS_HLSL==0 && !defined(NIS_GLSL) +#define NIS_GLSL 1 +#endif +#if NIS_HLSL_6_2 || (!NIS_GLSL && !NIS_HLSL) +#if defined(NIS_HLSL) +#undef NIS_HLSL +#endif +#define NIS_HLSL 1 +#endif +#if NIS_HLSL && NIS_GLSL +#undef NIS_GLSL +#define NIS_GLSL 0 +#endif + +// Half precision +#ifndef NIS_USE_HALF_PRECISION +#define NIS_USE_HALF_PRECISION 0 +#endif + +#if NIS_HLSL +// Generic type and function aliases for HLSL +#define NVF float +#define NVF2 float2 +#define NVF3 float3 +#define NVF4 float4 +#define NVI int +#define NVI2 int2 +#define NVU uint +#define NVU2 uint2 +#define NVB bool +#if NIS_USE_HALF_PRECISION +#if NIS_HLSL_6_2 +#define NVH float16_t +#define NVH2 float16_t2 +#define NVH3 float16_t3 +#define NVH4 float16_t4 +#else +#define NVH min16float +#define NVH2 min16float2 +#define NVH3 min16float3 +#define NVH4 min16float4 +#endif // NIS_HLSL_6_2 +#else // FP32 types +#define NVH NVF +#define NVH2 NVF2 +#define NVH3 NVF3 +#define NVH4 NVF4 +#endif // NIS_USE_HALF_PRECISION +#define NVSHARED groupshared +#define NVTEX_LOAD(x, pos) x[pos] +#define NVTEX_SAMPLE(x, sampler, pos) x.SampleLevel(sampler, pos, 0) +#define NVTEX_SAMPLE_RED(x, sampler, pos) x.GatherRed(sampler, pos) +#define NVTEX_SAMPLE_GREEN(x, sampler, pos) x.GatherGreen(sampler, pos) +#define NVTEX_SAMPLE_BLUE(x, sampler, pos) x.GatherBlue(sampler, pos) +#define NVTEX_STORE(x, pos, v) x[pos] = v +#ifndef NIS_UNROLL +#define NIS_UNROLL [unroll] +#endif +#endif // NIS_HLSL + +// Generic type and function aliases for GLSL +#if NIS_GLSL +#define NVF float +#define NVF2 vec2 +#define NVF3 vec3 +#define NVF4 vec4 +#define NVI int +#define NVI2 ivec2 +#define NVU uint +#define NVU2 uvec2 +#define NVB bool +#if NIS_USE_HALF_PRECISION +#define NVH float16_t +#define NVH2 f16vec2 +#define NVH3 f16vec3 +#define NVH4 f16vec4 +#else // FP32 types +#define NVH NVF +#define NVH2 NVF2 +#define NVH3 NVF3 +#define NVH4 NVF4 +#endif // NIS_USE_HALF_PRECISION +#define NVSHARED shared +#define NVTEX_LOAD(arr, pos) arr[(pos).y][(pos).x] // Ryujinx: coef from inline const arrays (param renamed to avoid .x swizzle collision) +#define NVTEX_SAMPLE(x, sampler, pos) textureLod(x, pos, 0.0) // Ryujinx: combined sampler2D +#define NVTEX_SAMPLE_RED(x, sampler, pos) textureGather(x, pos, 0) +#define NVTEX_SAMPLE_GREEN(x, sampler, pos) textureGather(x, pos, 1) +#define NVTEX_SAMPLE_BLUE(x, sampler, pos) textureGather(x, pos, 2) +#ifndef NIS_OUTPUT_TRANSFORM +#define NIS_OUTPUT_TRANSFORM(c) (c) // Ryujinx: HDR wrapper overrides this to tonemap to scRGB +#endif +#define NVTEX_STORE(x, pos, v) imageStore(x, NVI2(pos), NIS_OUTPUT_TRANSFORM(v)) +#define saturate(x) clamp(x, 0, 1) +#define lerp(a, b, x) mix(a, b, x) +#define GroupMemoryBarrierWithGroupSync() groupMemoryBarrier(); barrier() +#ifndef NIS_UNROLL +#define NIS_UNROLL +#endif +#endif // NIS_GLSL + +#ifndef NIS_UNROLL_INNER +#define NIS_UNROLL_INNER NIS_UNROLL +#endif + +// Texture gather +#ifndef NIS_TEXTURE_GATHER +#define NIS_TEXTURE_GATHER 0 +#endif + +// NIS Scaling +#define NIS_SCALE_INT 1 +#define NIS_SCALE_FLOAT NVF(1.f) + +// NIS output clamp +#if NIS_CLAMP_OUTPUT +#if NIS_HDR_MODE == NIS_HDR_MODE_LINEAR +#define NVCLAMP(x) ( clamp(x, 0.0f, 12.5f) ) +#else +#define NVCLAMP(x) ( saturate(x) ) +#endif +#else +#define NVCLAMP(x) (x) +#endif + +NVF getY(NVF3 rgba) +{ +#if NIS_HDR_MODE == NIS_HDR_MODE_PQ + return NVF(0.262f) * rgba.x + NVF(0.678f) * rgba.y + NVF(0.0593f) * rgba.z; +#elif NIS_HDR_MODE == NIS_HDR_MODE_LINEAR + return sqrt(NVF(0.2126f) * rgba.x + NVF(0.7152f) * rgba.y + NVF(0.0722f) * rgba.z) * kHDRCompressionFactor; +#else + return NVF(0.2126f) * rgba.x + NVF(0.7152f) * rgba.y + NVF(0.0722f) * rgba.z; +#endif +} + +NVF getYLinear(NVF3 rgba) +{ + return NVF(0.2126f) * rgba.x + NVF(0.7152f) * rgba.y + NVF(0.0722f) * rgba.z; +} + +NVF3 YUVtoRGB(NVF3 yuv) +{ + float y = yuv.x - 16.0f / 255.0f; + float u = yuv.y - 128.0f / 255.0f; + float v = yuv.z - 128.0f / 255.0f; + NVF3 rgb; + rgb.x = saturate(1.164f * y + 1.596f * v); + rgb.y = saturate(1.164f * y - 0.392f * u - 0.813f * v); + rgb.z = saturate(1.164f * y + 2.017f * u); + return rgb; +} + +#if NIS_SCALER +NVF4 GetEdgeMap(NVF p[4][4], NVI i, NVI j) +#else +NVF4 GetEdgeMap(NVF p[5][5], NVI i, NVI j) +#endif +{ + const NVF g_0 = abs(p[0 + i][0 + j] + p[0 + i][1 + j] + p[0 + i][2 + j] - p[2 + i][0 + j] - p[2 + i][1 + j] - p[2 + i][2 + j]); + const NVF g_45 = abs(p[1 + i][0 + j] + p[0 + i][0 + j] + p[0 + i][1 + j] - p[2 + i][1 + j] - p[2 + i][2 + j] - p[1 + i][2 + j]); + const NVF g_90 = abs(p[0 + i][0 + j] + p[1 + i][0 + j] + p[2 + i][0 + j] - p[0 + i][2 + j] - p[1 + i][2 + j] - p[2 + i][2 + j]); + const NVF g_135 = abs(p[1 + i][0 + j] + p[2 + i][0 + j] + p[2 + i][1 + j] - p[0 + i][1 + j] - p[0 + i][2 + j] - p[1 + i][2 + j]); + + const NVF g_0_90_max = max(g_0, g_90); + const NVF g_0_90_min = min(g_0, g_90); + const NVF g_45_135_max = max(g_45, g_135); + const NVF g_45_135_min = min(g_45, g_135); + + NVF e_0_90 = 0; + NVF e_45_135 = 0; + + if (g_0_90_max + g_45_135_max == 0) + { + return NVF4(0, 0, 0, 0); + } + + e_0_90 = min(g_0_90_max / (g_0_90_max + g_45_135_max), 1.0f); + e_45_135 = 1.0f - e_0_90; + + NVB c_0_90 = (g_0_90_max > (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 000000000..b296cd6d8 Binary files /dev/null and b/src/Ryujinx.Graphics.Vulkan/Effects/Shaders/NisScaling.spv differ 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 000000000..e9567dedb Binary files /dev/null and b/src/Ryujinx.Graphics.Vulkan/Effects/Shaders/NisScalingHdr.spv differ diff --git a/src/Ryujinx.Graphics.Vulkan/Effects/Shaders/nis_coef.glsl b/src/Ryujinx.Graphics.Vulkan/Effects/Shaders/nis_coef.glsl new file mode 100644 index 000000000..f33ed3b9d --- /dev/null +++ b/src/Ryujinx.Graphics.Vulkan/Effects/Shaders/nis_coef.glsl @@ -0,0 +1,157 @@ +// 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. +// +// NIS coefficient tables, extracted verbatim from NVIDIA NIS_Config.h (v1.0.3, MIT). +// Layout: [phase 0..63][col 0=taps0-3, 1=taps4-7] matching texelFetch(coef, ivec2(vIdx,phase)). + +const vec4 coef_scaler[64][2] = vec4[64][2]( + vec4[2](vec4(0.0, 0.0, 1.0, 0.0), vec4(0.0, 0.0, 0.0, 0.0)), + vec4[2](vec4(0.0029, -0.0127, 1.0, 0.0132), vec4(-0.0034, 0.0, 0.0, 0.0)), + vec4[2](vec4(0.0063, -0.0249, 0.9985, 0.0269), vec4(-0.0068, 0.0, 0.0, 0.0)), + vec4[2](vec4(0.0088, -0.0361, 0.9956, 0.0415), vec4(-0.0103, 0.0005, 0.0, 0.0)), + vec4[2](vec4(0.0117, -0.0474, 0.9932, 0.0562), vec4(-0.0142, 0.0005, 0.0, 0.0)), + vec4[2](vec4(0.0142, -0.0576, 0.9897, 0.0713), vec4(-0.0181, 0.0005, 0.0, 0.0)), + vec4[2](vec4(0.0166, -0.0674, 0.9844, 0.0874), vec4(-0.022, 0.001, 0.0, 0.0)), + vec4[2](vec4(0.0186, -0.0762, 0.9785, 0.104), vec4(-0.0264, 0.0015, 0.0, 0.0)), + vec4[2](vec4(0.0205, -0.085, 0.9727, 0.1206), vec4(-0.0308, 0.002, 0.0, 0.0)), + vec4[2](vec4(0.0225, -0.0928, 0.9648, 0.1382), vec4(-0.0352, 0.0024, 0.0, 0.0)), + vec4[2](vec4(0.0239, -0.1006, 0.9575, 0.1558), vec4(-0.0396, 0.0029, 0.0, 0.0)), + vec4[2](vec4(0.0254, -0.1074, 0.9487, 0.1738), vec4(-0.0439, 0.0034, 0.0, 0.0)), + vec4[2](vec4(0.0264, -0.1138, 0.939, 0.1929), vec4(-0.0488, 0.0044, 0.0, 0.0)), + vec4[2](vec4(0.0278, -0.1191, 0.9282, 0.2119), vec4(-0.0537, 0.0049, 0.0, 0.0)), + vec4[2](vec4(0.0288, -0.1245, 0.917, 0.231), vec4(-0.0581, 0.0059, 0.0, 0.0)), + vec4[2](vec4(0.0293, -0.1294, 0.9058, 0.251), vec4(-0.063, 0.0063, 0.0, 0.0)), + vec4[2](vec4(0.0303, -0.1333, 0.8926, 0.271), vec4(-0.0679, 0.0073, 0.0, 0.0)), + vec4[2](vec4(0.0308, -0.1367, 0.8789, 0.2915), vec4(-0.0728, 0.0083, 0.0, 0.0)), + vec4[2](vec4(0.0308, -0.1401, 0.8657, 0.312), vec4(-0.0776, 0.0093, 0.0, 0.0)), + vec4[2](vec4(0.0313, -0.1426, 0.8506, 0.333), vec4(-0.0825, 0.0103, 0.0, 0.0)), + vec4[2](vec4(0.0313, -0.1445, 0.8354, 0.354), vec4(-0.0874, 0.0112, 0.0, 0.0)), + vec4[2](vec4(0.0313, -0.146, 0.8193, 0.3755), vec4(-0.0923, 0.0122, 0.0, 0.0)), + vec4[2](vec4(0.0313, -0.147, 0.8022, 0.3965), vec4(-0.0967, 0.0137, 0.0, 0.0)), + vec4[2](vec4(0.0308, -0.1479, 0.7856, 0.4185), vec4(-0.1016, 0.0146, 0.0, 0.0)), + vec4[2](vec4(0.0303, -0.1479, 0.7681, 0.4399), vec4(-0.106, 0.0156, 0.0, 0.0)), + vec4[2](vec4(0.0298, -0.1479, 0.7505, 0.4614), vec4(-0.1104, 0.0166, 0.0, 0.0)), + vec4[2](vec4(0.0293, -0.147, 0.7314, 0.4829), vec4(-0.1147, 0.0181, 0.0, 0.0)), + vec4[2](vec4(0.0288, -0.146, 0.7119, 0.5049), vec4(-0.1187, 0.019, 0.0, 0.0)), + vec4[2](vec4(0.0278, -0.1445, 0.6929, 0.5264), vec4(-0.1226, 0.02, 0.0, 0.0)), + vec4[2](vec4(0.0273, -0.1431, 0.6724, 0.5479), vec4(-0.126, 0.0215, 0.0, 0.0)), + vec4[2](vec4(0.0264, -0.1411, 0.6528, 0.5693), vec4(-0.1299, 0.0225, 0.0, 0.0)), + vec4[2](vec4(0.0254, -0.1387, 0.6323, 0.5903), vec4(-0.1328, 0.0234, 0.0, 0.0)), + vec4[2](vec4(0.0244, -0.1357, 0.6113, 0.6113), vec4(-0.1357, 0.0244, 0.0, 0.0)), + vec4[2](vec4(0.0234, -0.1328, 0.5903, 0.6323), vec4(-0.1387, 0.0254, 0.0, 0.0)), + vec4[2](vec4(0.0225, -0.1299, 0.5693, 0.6528), vec4(-0.1411, 0.0264, 0.0, 0.0)), + vec4[2](vec4(0.0215, -0.126, 0.5479, 0.6724), vec4(-0.1431, 0.0273, 0.0, 0.0)), + vec4[2](vec4(0.02, -0.1226, 0.5264, 0.6929), vec4(-0.1445, 0.0278, 0.0, 0.0)), + vec4[2](vec4(0.019, -0.1187, 0.5049, 0.7119), vec4(-0.146, 0.0288, 0.0, 0.0)), + vec4[2](vec4(0.0181, -0.1147, 0.4829, 0.7314), vec4(-0.147, 0.0293, 0.0, 0.0)), + vec4[2](vec4(0.0166, -0.1104, 0.4614, 0.7505), vec4(-0.1479, 0.0298, 0.0, 0.0)), + vec4[2](vec4(0.0156, -0.106, 0.4399, 0.7681), vec4(-0.1479, 0.0303, 0.0, 0.0)), + vec4[2](vec4(0.0146, -0.1016, 0.4185, 0.7856), vec4(-0.1479, 0.0308, 0.0, 0.0)), + vec4[2](vec4(0.0137, -0.0967, 0.3965, 0.8022), vec4(-0.147, 0.0313, 0.0, 0.0)), + vec4[2](vec4(0.0122, -0.0923, 0.3755, 0.8193), vec4(-0.146, 0.0313, 0.0, 0.0)), + vec4[2](vec4(0.0112, -0.0874, 0.354, 0.8354), vec4(-0.1445, 0.0313, 0.0, 0.0)), + vec4[2](vec4(0.0103, -0.0825, 0.333, 0.8506), vec4(-0.1426, 0.0313, 0.0, 0.0)), + vec4[2](vec4(0.0093, -0.0776, 0.312, 0.8657), vec4(-0.1401, 0.0308, 0.0, 0.0)), + vec4[2](vec4(0.0083, -0.0728, 0.2915, 0.8789), vec4(-0.1367, 0.0308, 0.0, 0.0)), + vec4[2](vec4(0.0073, -0.0679, 0.271, 0.8926), vec4(-0.1333, 0.0303, 0.0, 0.0)), + vec4[2](vec4(0.0063, -0.063, 0.251, 0.9058), vec4(-0.1294, 0.0293, 0.0, 0.0)), + vec4[2](vec4(0.0059, -0.0581, 0.231, 0.917), vec4(-0.1245, 0.0288, 0.0, 0.0)), + vec4[2](vec4(0.0049, -0.0537, 0.2119, 0.9282), vec4(-0.1191, 0.0278, 0.0, 0.0)), + vec4[2](vec4(0.0044, -0.0488, 0.1929, 0.939), vec4(-0.1138, 0.0264, 0.0, 0.0)), + vec4[2](vec4(0.0034, -0.0439, 0.1738, 0.9487), vec4(-0.1074, 0.0254, 0.0, 0.0)), + vec4[2](vec4(0.0029, -0.0396, 0.1558, 0.9575), vec4(-0.1006, 0.0239, 0.0, 0.0)), + vec4[2](vec4(0.0024, -0.0352, 0.1382, 0.9648), vec4(-0.0928, 0.0225, 0.0, 0.0)), + vec4[2](vec4(0.002, -0.0308, 0.1206, 0.9727), vec4(-0.085, 0.0205, 0.0, 0.0)), + vec4[2](vec4(0.0015, -0.0264, 0.104, 0.9785), vec4(-0.0762, 0.0186, 0.0, 0.0)), + vec4[2](vec4(0.001, -0.022, 0.0874, 0.9844), vec4(-0.0674, 0.0166, 0.0, 0.0)), + vec4[2](vec4(0.0005, -0.0181, 0.0713, 0.9897), vec4(-0.0576, 0.0142, 0.0, 0.0)), + vec4[2](vec4(0.0005, -0.0142, 0.0562, 0.9932), vec4(-0.0474, 0.0117, 0.0, 0.0)), + vec4[2](vec4(0.0005, -0.0103, 0.0415, 0.9956), vec4(-0.0361, 0.0088, 0.0, 0.0)), + vec4[2](vec4(0.0, -0.0068, 0.0269, 0.9985), vec4(-0.0249, 0.0063, 0.0, 0.0)), + vec4[2](vec4(0.0, -0.0034, 0.0132, 1.0), vec4(-0.0127, 0.0029, 0.0, 0.0)) +); + +const vec4 coef_usm[64][2] = vec4[64][2]( + vec4[2](vec4(0.0, -0.6001, 1.2002, -0.6001), vec4(0.0, 0.0, 0.0, 0.0)), + vec4[2](vec4(0.0029, -0.6084, 1.1987, -0.5903), vec4(-0.0029, 0.0, 0.0, 0.0)), + vec4[2](vec4(0.0049, -0.6147, 1.1958, -0.5791), vec4(-0.0068, 0.0005, 0.0, 0.0)), + vec4[2](vec4(0.0073, -0.6196, 1.189, -0.5659), vec4(-0.0103, 0.0, 0.0, 0.0)), + vec4[2](vec4(0.0093, -0.6235, 1.1802, -0.5513), vec4(-0.0151, 0.0, 0.0, 0.0)), + vec4[2](vec4(0.0112, -0.6265, 1.1699, -0.5352), vec4(-0.0195, 0.0005, 0.0, 0.0)), + vec4[2](vec4(0.0122, -0.627, 1.1582, -0.5181), vec4(-0.0259, 0.0005, 0.0, 0.0)), + vec4[2](vec4(0.0142, -0.6284, 1.1455, -0.5005), vec4(-0.0317, 0.0005, 0.0, 0.0)), + vec4[2](vec4(0.0156, -0.6265, 1.1274, -0.479), vec4(-0.0386, 0.0005, 0.0, 0.0)), + vec4[2](vec4(0.0166, -0.6235, 1.1089, -0.457), vec4(-0.0454, 0.001, 0.0, 0.0)), + vec4[2](vec4(0.0176, -0.6187, 1.0879, -0.4346), vec4(-0.0532, 0.001, 0.0, 0.0)), + vec4[2](vec4(0.0181, -0.6138, 1.0659, -0.4102), vec4(-0.0615, 0.0015, 0.0, 0.0)), + vec4[2](vec4(0.019, -0.6069, 1.0405, -0.3843), vec4(-0.0698, 0.0015, 0.0, 0.0)), + vec4[2](vec4(0.0195, -0.6006, 1.0161, -0.3574), vec4(-0.0796, 0.002, 0.0, 0.0)), + vec4[2](vec4(0.02, -0.5928, 0.9893, -0.3286), vec4(-0.0898, 0.0024, 0.0, 0.0)), + vec4[2](vec4(0.02, -0.582, 0.958, -0.2988), vec4(-0.1001, 0.0029, 0.0, 0.0)), + vec4[2](vec4(0.02, -0.5728, 0.9292, -0.269), vec4(-0.1104, 0.0034, 0.0, 0.0)), + vec4[2](vec4(0.02, -0.562, 0.8975, -0.2368), vec4(-0.1226, 0.0039, 0.0, 0.0)), + vec4[2](vec4(0.0205, -0.5498, 0.8643, -0.2046), vec4(-0.1343, 0.0044, 0.0, 0.0)), + vec4[2](vec4(0.02, -0.5371, 0.8301, -0.1709), vec4(-0.1465, 0.0049, 0.0, 0.0)), + vec4[2](vec4(0.0195, -0.5239, 0.7944, -0.1367), vec4(-0.1587, 0.0054, 0.0, 0.0)), + vec4[2](vec4(0.0195, -0.5107, 0.7598, -0.1021), vec4(-0.1724, 0.0059, 0.0, 0.0)), + vec4[2](vec4(0.019, -0.4966, 0.7231, -0.0649), vec4(-0.1865, 0.0063, 0.0, 0.0)), + vec4[2](vec4(0.0186, -0.4819, 0.6846, -0.0288), vec4(-0.1997, 0.0068, 0.0, 0.0)), + vec4[2](vec4(0.0186, -0.4668, 0.646, 0.0093), vec4(-0.2144, 0.0073, 0.0, 0.0)), + vec4[2](vec4(0.0176, -0.4507, 0.6055, 0.0479), vec4(-0.229, 0.0083, 0.0, 0.0)), + vec4[2](vec4(0.0171, -0.437, 0.5693, 0.0859), vec4(-0.2446, 0.0088, 0.0, 0.0)), + vec4[2](vec4(0.0161, -0.4199, 0.5283, 0.1255), vec4(-0.2598, 0.0098, 0.0, 0.0)), + vec4[2](vec4(0.0161, -0.4048, 0.4883, 0.1655), vec4(-0.2754, 0.0103, 0.0, 0.0)), + vec4[2](vec4(0.0151, -0.3887, 0.4497, 0.2041), vec4(-0.291, 0.0107, 0.0, 0.0)), + vec4[2](vec4(0.0142, -0.3711, 0.4072, 0.2446), vec4(-0.3066, 0.0117, 0.0, 0.0)), + vec4[2](vec4(0.0137, -0.3555, 0.3672, 0.2852), vec4(-0.3228, 0.0122, 0.0, 0.0)), + vec4[2](vec4(0.0132, -0.3394, 0.3262, 0.3262), vec4(-0.3394, 0.0132, 0.0, 0.0)), + vec4[2](vec4(0.0122, -0.3228, 0.2852, 0.3672), vec4(-0.3555, 0.0137, 0.0, 0.0)), + vec4[2](vec4(0.0117, -0.3066, 0.2446, 0.4072), vec4(-0.3711, 0.0142, 0.0, 0.0)), + vec4[2](vec4(0.0107, -0.291, 0.2041, 0.4497), vec4(-0.3887, 0.0151, 0.0, 0.0)), + vec4[2](vec4(0.0103, -0.2754, 0.1655, 0.4883), vec4(-0.4048, 0.0161, 0.0, 0.0)), + vec4[2](vec4(0.0098, -0.2598, 0.1255, 0.5283), vec4(-0.4199, 0.0161, 0.0, 0.0)), + vec4[2](vec4(0.0088, -0.2446, 0.0859, 0.5693), vec4(-0.437, 0.0171, 0.0, 0.0)), + vec4[2](vec4(0.0083, -0.229, 0.0479, 0.6055), vec4(-0.4507, 0.0176, 0.0, 0.0)), + vec4[2](vec4(0.0073, -0.2144, 0.0093, 0.646), vec4(-0.4668, 0.0186, 0.0, 0.0)), + vec4[2](vec4(0.0068, -0.1997, -0.0288, 0.6846), vec4(-0.4819, 0.0186, 0.0, 0.0)), + vec4[2](vec4(0.0063, -0.1865, -0.0649, 0.7231), vec4(-0.4966, 0.019, 0.0, 0.0)), + vec4[2](vec4(0.0059, -0.1724, -0.1021, 0.7598), vec4(-0.5107, 0.0195, 0.0, 0.0)), + vec4[2](vec4(0.0054, -0.1587, -0.1367, 0.7944), vec4(-0.5239, 0.0195, 0.0, 0.0)), + vec4[2](vec4(0.0049, -0.1465, -0.1709, 0.8301), vec4(-0.5371, 0.02, 0.0, 0.0)), + vec4[2](vec4(0.0044, -0.1343, -0.2046, 0.8643), vec4(-0.5498, 0.0205, 0.0, 0.0)), + vec4[2](vec4(0.0039, -0.1226, -0.2368, 0.8975), vec4(-0.562, 0.02, 0.0, 0.0)), + vec4[2](vec4(0.0034, -0.1104, -0.269, 0.9292), vec4(-0.5728, 0.02, 0.0, 0.0)), + vec4[2](vec4(0.0029, -0.1001, -0.2988, 0.958), vec4(-0.582, 0.02, 0.0, 0.0)), + vec4[2](vec4(0.0024, -0.0898, -0.3286, 0.9893), vec4(-0.5928, 0.02, 0.0, 0.0)), + vec4[2](vec4(0.002, -0.0796, -0.3574, 1.0161), vec4(-0.6006, 0.0195, 0.0, 0.0)), + vec4[2](vec4(0.0015, -0.0698, -0.3843, 1.0405), vec4(-0.6069, 0.019, 0.0, 0.0)), + vec4[2](vec4(0.0015, -0.0615, -0.4102, 1.0659), vec4(-0.6138, 0.0181, 0.0, 0.0)), + vec4[2](vec4(0.001, -0.0532, -0.4346, 1.0879), vec4(-0.6187, 0.0176, 0.0, 0.0)), + vec4[2](vec4(0.001, -0.0454, -0.457, 1.1089), vec4(-0.6235, 0.0166, 0.0, 0.0)), + vec4[2](vec4(0.0005, -0.0386, -0.479, 1.1274), vec4(-0.6265, 0.0156, 0.0, 0.0)), + vec4[2](vec4(0.0005, -0.0317, -0.5005, 1.1455), vec4(-0.6284, 0.0142, 0.0, 0.0)), + vec4[2](vec4(0.0005, -0.0259, -0.5181, 1.1582), vec4(-0.627, 0.0122, 0.0, 0.0)), + vec4[2](vec4(0.0005, -0.0195, -0.5352, 1.1699), vec4(-0.6265, 0.0112, 0.0, 0.0)), + vec4[2](vec4(0.0, -0.0151, -0.5513, 1.1802), vec4(-0.6235, 0.0093, 0.0, 0.0)), + vec4[2](vec4(0.0, -0.0103, -0.5659, 1.189), vec4(-0.6196, 0.0073, 0.0, 0.0)), + vec4[2](vec4(0.0005, -0.0068, -0.5791, 1.1958), vec4(-0.6147, 0.0049, 0.0, 0.0)), + vec4[2](vec4(0.0, -0.0029, -0.5903, 1.1987), vec4(-0.6084, 0.0029, 0.0, 0.0)) +); diff --git a/src/Ryujinx.Graphics.Vulkan/HelperShader.cs b/src/Ryujinx.Graphics.Vulkan/HelperShader.cs index 88e720cde..3f9ae9f10 100644 --- a/src/Ryujinx.Graphics.Vulkan/HelperShader.cs +++ b/src/Ryujinx.Graphics.Vulkan/HelperShader.cs @@ -369,7 +369,8 @@ namespace Ryujinx.Graphics.Vulkan 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); @@ -412,6 +413,7 @@ namespace Ryujinx.Graphics.Vulkan hdrParams[2] = curve; hdrParams[3] = gamma; hdrParams[4] = blend; + hdrParams[5] = whiten; hdrBuffer.Holder.SetDataUnchecked(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 b86284304..ff6bc4fdc 100644 Binary files a/src/Ryujinx.Graphics.Vulkan/Shaders/SpirvBinaries/ColorBlitHdrFragment.spv and b/src/Ryujinx.Graphics.Vulkan/Shaders/SpirvBinaries/ColorBlitHdrFragment.spv differ 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}"/> + + + + + +