NVIDIA Image Scaling v1.0.2 for mpv

NVScaler.glsl

// 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 v1.0.2 by NVIDIA
// ported to mpv by agyild

// Changelog
// Made it directly operate on LUMA plane, since the original shader was operating
// on LUMA by deriving it from RGB. This should cause a major increase in performance,
// especially on OpenGL 4.0+ renderers

//!HOOK LUMA
//!BIND HOOKED
//!BIND coef_scaler
//!BIND coef_usm
//!DESC NVIDIA Image Scaling and Sharpening v1.0.2
//!COMPUTE 32 24 256 1
//!WHEN OUTPUT.w OUTPUT.h * LUMA.w LUMA.h * / 1.0 >
//!WIDTH OUTPUT.w
//!HEIGHT OUTPUT.h

// User variables
#define SHARPNESS 0.25 // Amount of sharpening. 0.0 to 1.0.
#define NIS_THREAD_GROUP_SIZE 256 // May be set to 128 for better performance on NVIDIA hardware, otherwise set to 256. Don't forget to modify the COMPUTE directive accordingly as well (e.g., COMPUTE 32 24 128 1).
#define NIS_HDR_MODE 0 // Must be set to 1 for content with PQ colorspace. 0 or 1.

// Constant variables
#define NIS_BLOCK_WIDTH 32
#define NIS_BLOCK_HEIGHT 24
#define kPhaseCount  64
#define kFilterSize  6
#define kSupportSize 6
#define kPadSize     kSupportSize
#define NIS_SCALE_INT 1
#define NIS_SCALE_FLOAT 1.0f
#define kTilePitch (NIS_BLOCK_WIDTH + kPadSize)
#define kTileSize (kTilePitch * (NIS_BLOCK_HEIGHT + kPadSize))
#define kEdgeMapPitch (NIS_BLOCK_WIDTH + 2)
#define kEdgeMapSize (kEdgeMapPitch * (NIS_BLOCK_HEIGHT + 2))
const float sharpen_slider = clamp(SHARPNESS, 0.0f, 1.0f) - 0.5f;
const float MaxScale = (sharpen_slider >= 0.0f) ? 1.25f : 1.75f;
const float MinScale = (sharpen_slider >= 0.0f) ? 1.25f : 1.0f;
const float LimitScale = (sharpen_slider >= 0.0f) ? 1.25f : 1.0f;
const float kDetectRatio = 2 * 1127.f / 1024.f;
const float kDetectThres = (bool(NIS_HDR_MODE) ? 32.0f : 64.0f) / 1024.0f;
const float kMinContrastRatio = bool(NIS_HDR_MODE) ? 1.5f : 2.0f;
const float kMaxContrastRatio = bool(NIS_HDR_MODE) ? 5.0f : 10.0f;
const float kSharpStartY = bool(NIS_HDR_MODE) ? 0.35f : 0.45f;
const float kSharpEndY = bool(NIS_HDR_MODE) ? 0.55f : 0.9f;
const float kSharpStrengthMin = max(0.0f, 0.4f + sharpen_slider * MinScale * (bool(NIS_HDR_MODE) ? 1.1f : 1.2));
const float kSharpStrengthMax = ((bool(NIS_HDR_MODE) ? 2.2f : 1.6f) + sharpen_slider * MaxScale * 1.8f);
const float kSharpLimitMin = max((bool(NIS_HDR_MODE) ? 0.06f :0.1f), (bool(NIS_HDR_MODE) ? 0.1f : 0.14f) + sharpen_slider * LimitScale * (bool(NIS_HDR_MODE) ? 0.28f : 0.32f));
const float kSharpLimitMax = ((bool(NIS_HDR_MODE) ? 0.6f : 0.5f) + sharpen_slider * LimitScale * 0.6f);
const float kRatioNorm = 1.0f / (kMaxContrastRatio - kMinContrastRatio);
const float kSharpScaleY = 1.0f / (kSharpEndY - kSharpStartY);
const float kSharpStrengthScale = kSharpStrengthMax - kSharpStrengthMin;
const float kSharpLimitScale = kSharpLimitMax - kSharpLimitMin;
const float kContrastBoost = 1.0f;
const float kEps = 1.0f;
#define kScaleX (HOOKED_size.x / target_size.x)
#define kScaleY (HOOKED_size.y / target_size.y)
#define kSrcNormX HOOKED_pt.x
#define kSrcNormY HOOKED_pt.y

// HLSL to GLSL macros
#define saturate(x) clamp(x, 0, 1)
#define lerp(a, b, x) mix(a, b, x)

// CS Shared variables
shared float shPixelsY[kTileSize];
shared float shCoefScaler[kPhaseCount][kFilterSize];
shared float shCoefUSM[kPhaseCount][kFilterSize];
shared vec4 shEdgeMap[kEdgeMapSize];

// Shader code

vec4 GetEdgeMap(float p[4][4], int i, int j) {
	const float 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 float 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 float 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 float 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 float g_0_90_max = max(g_0, g_90);
	const float g_0_90_min = min(g_0, g_90);
	const float g_45_135_max = max(g_45, g_135);
	const float g_45_135_min = min(g_45, g_135);

	float e_0_90 = 0;
	float e_45_135 = 0;

	if (g_0_90_max + g_45_135_max == 0)
	{
		return vec4(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;

	bool 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);
	bool 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);
	bool c_g_0_90 = g_0_90_max == g_0;
	bool c_g_45_135 = g_45_135_max == g_45;

	float f_e_0_90 = (c_0_90 && c_45_135) ? e_0_90 : 1.0f;
	float f_e_45_135 = (c_0_90 && c_45_135) ? e_45_135 : 1.0f;

	float weight_0 = (c_0_90 && c_g_0_90) ? f_e_0_90 : 0.0f;
	float weight_90 = (c_0_90 && !c_g_0_90) ? f_e_0_90 : 0.0f;
	float weight_45 = (c_45_135 && c_g_45_135) ? f_e_45_135 : 0.0f;
	float weight_135 = (c_45_135 && !c_g_45_135) ? f_e_45_135 : 0.0f;

	return vec4(weight_0, weight_90, weight_45, weight_135);
}

void LoadFilterBanksSh(int i0, int di) {
	// Load up filter banks to shared memory
	// The work is spread over (kPhaseCount * 2) threads
	for (int i = i0; i < kPhaseCount * 2; i += di)
	{
		int phase = i >> 1;
		int vIdx = i & 1;

		vec4 v = vec4(texelFetch(coef_scaler, ivec2(vIdx, phase), 0));
		int 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 = vec4(texelFetch(coef_usm, ivec2(vIdx, phase), 0));
		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;
		}
	}
}

float CalcLTI(float p0, float p1, float p2, float p3, float p4, float p5, int phase_index)
{
	const bool selector = (phase_index <= kPhaseCount / 2);
	float sel = selector ? p0 : p3;
	const float a_min = min(min(p1, p2), sel);
	const float a_max = max(max(p1, p2), sel);
	sel = selector ? p2 : p5;
	const float b_min = min(min(p3, p4), sel);
	const float b_max = max(max(p3, p4), sel);

	const float a_cont = a_max - a_min;
	const float b_cont = b_max - b_min;

	const float cont_ratio = max(a_cont, b_cont) / (min(a_cont, b_cont) + kEps);
	return (1.0f - saturate((cont_ratio - kMinContrastRatio) * kRatioNorm)) * kContrastBoost;
}

vec4 GetInterpEdgeMap(const vec4 edge[2][2], float phase_frac_x, float phase_frac_y)
{
	vec4 h0 = lerp(edge[0][0], edge[0][1], phase_frac_x);
	vec4 h1 = lerp(edge[1][0], edge[1][1], phase_frac_x);
	return lerp(h0, h1, phase_frac_y);
}

float EvalPoly6(const float pxl[6], int phase_int)
{
	float y = 0.f;
	{
		for (int i = 0; i < 6; ++i)
		{
			y += shCoefScaler[phase_int][i] * pxl[i];
		}
	}
	float y_usm = 0.f;
	{
		for (int i = 0; i < 6; ++i)
		{
			y_usm += shCoefUSM[phase_int][i] * pxl[i];
		}
	}

	// let's compute a piece-wise ramp based on luma
	const float y_scale = 1.0f - saturate((y * (1.0f / NIS_SCALE_FLOAT) - kSharpStartY) * kSharpScaleY);

	// scale the ramp to sharpen as a function of luma
	const float y_sharpness = y_scale * kSharpStrengthScale + kSharpStrengthMin;

	y_usm *= y_sharpness;

	// scale the ramp to limit USM as a function of luma
	const float 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;
}

float FilterNormal(const float p[6][6], int phase_x_frac_int, int phase_y_frac_int)
{
	float h_acc = 0.0f;
	for (int j = 0; j < 6; ++j)
	{
		float v_acc = 0.0f;
		for (int 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;
}

float AddDirFilters(float p[6][6], float phase_x_frac, float phase_y_frac, int phase_x_frac_int, int phase_y_frac_int, vec4 w)
{
	float f = 0.f;
	if (w.x > 0.0f)
	{
		// 0 deg filter
		float interp0Deg[6];
		{
			for (int 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
		float interp90Deg[6];
		{
			for (int 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
		float pphase_b45;
		pphase_b45 = 0.5f + 0.5f * (phase_x_frac - phase_y_frac);

		float 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;
			float a = (pphase_b45 >= 0.f) ? p[0][2] : p[2][0];
			float b = (pphase_b45 >= 0.f) ? p[1][3] : p[3][1];
			float c = (pphase_b45 >= 0.f) ? p[2][4] : p[4][2];
			float 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));
		}

		float interp45Deg[6];
		float pphase_p45 = phase_x_frac + phase_y_frac;
		if (pphase_p45 >= 1)
		{
			for (int i = 0; i < 6; i++)
			{
				interp45Deg[i] = temp_interp45Deg[i + 1];
			}
			pphase_p45 = pphase_p45 - 1;
		}
		else
		{
			for (int i = 0; i < 6; i++)
			{
				interp45Deg[i] = temp_interp45Deg[i];
			}
		}
		f += EvalPoly6(interp45Deg, int(pphase_p45 * 64)) * w.z;
	}
	
	if (w.w > 0.0f)
	{
		//135 deg filter
		float pphase_b135 = 0.5f * (phase_x_frac + phase_y_frac);

		float 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;
			float a = (pphase_b135 >= 0.f) ? p[5][2] : p[3][0];
			float b = (pphase_b135 >= 0.f) ? p[4][3] : p[2][1];
			float c = (pphase_b135 >= 0.f) ? p[3][4] : p[1][2];
			float 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));
		}


		float interp135Deg[6];
		float pphase_p135 = 1 + (phase_x_frac - phase_y_frac);
		if (pphase_p135 >= 1)
		{
			for (int i = 0; i < 6; ++i)
			{
				interp135Deg[i] = temp_interp135Deg[i + 1];
			}
			pphase_p135 = pphase_p135 - 1;
		}
		else
		{
			for (int i = 0; i < 6; ++i)
			{
				interp135Deg[i] = temp_interp135Deg[i];
			}
		}
		f += EvalPoly6(interp135Deg, int(pphase_p135 * 64)) * w.w;
	}
	return f;
}

void hook()
{
	uvec2 blockIdx = gl_WorkGroupID.xy;
	uint threadIdx = gl_LocalInvocationID.x;

	// Figure out the range of pixels from input image that would be needed to be loaded for this thread-block
	int dstBlockX = int(NIS_BLOCK_WIDTH * blockIdx.x);
	int dstBlockY = int(NIS_BLOCK_HEIGHT * blockIdx.y);

	const int srcBlockStartX = int(floor((dstBlockX + 0.5f) * kScaleX - 0.5f));
	const int srcBlockStartY = int(floor((dstBlockY + 0.5f) * kScaleY - 0.5f));
	const int srcBlockEndX = int(ceil((dstBlockX + NIS_BLOCK_WIDTH + 0.5f) * kScaleX - 0.5f));
	const int srcBlockEndY = int(ceil((dstBlockY + NIS_BLOCK_HEIGHT + 0.5f) * kScaleY - 0.5f));

	int numTilePixelsX = srcBlockEndX - srcBlockStartX + kSupportSize - 1;
	int 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 int numTilePixels = numTilePixelsX * numTilePixelsY;

	// calculate the equivalent values for the edge map
	const int numEdgeMapPixelsX = numTilePixelsX - kSupportSize + 2;
	const int numEdgeMapPixelsY = numTilePixelsY - kSupportSize + 2;
	const int 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 (uint i = threadIdx * 2; i < uint(numTilePixels) >> 1; i += NIS_THREAD_GROUP_SIZE * 2)
		{
			uint py = (i / numTilePixelsX) * 2;
			uint px = i % numTilePixelsX;

			// 0.5 to be in the center of texel
			// - (kSupportSize - 1) / 2 to shift by the kernel support size
			float kShift = 0.5f - (kSupportSize - 1) / 2;

			const float tx = (srcBlockStartX + px + kShift) * kSrcNormX;
			const float ty = (srcBlockStartY + py + kShift) * kSrcNormY;

			float p[2][2];
#ifdef HOOKED_gather
			{
				const vec4 sY = HOOKED_gather(vec2(tx, ty), 0);

				p[0][0] = sY.w;
				p[0][1] = sY.z;
				p[1][0] = sY.x;
				p[1][1] = sY.y;
			}
#else
			for (int j = 0; j < 2; j++)
			{
				for (int k = 0; k < 2; k++)
				{
					const float px = HOOKED_tex(vec2(tx + k * kSrcNormX, ty + j * kSrcNormY)).r;
					p[j][k] = px;
				}
			}
#endif
			const uint idx = py * kTilePitch + px;
			shPixelsY[idx] = float(p[0][0]);
			shPixelsY[idx + 1] = float(p[0][1]);
			shPixelsY[idx + kTilePitch] = float(p[1][0]);
			shPixelsY[idx + kTilePitch + 1] = float(p[1][1]);
		}
	}

	groupMemoryBarrier();
	barrier();

	{
		// fill in the edge map of 2x2 pixels
		for (uint i = threadIdx * 2; i < uint(numEdgeMapPixels) >> 1; i += NIS_THREAD_GROUP_SIZE * 2)
		{
			uint py = (i / numEdgeMapPixelsX) * 2;
			uint px = i % numEdgeMapPixelsX;

			const uint edgeMapIdx = py * kEdgeMapPitch + px;

			uint tileCornerIdx = (py+1) * kTilePitch + px + 1;
			float p[4][4];
			for (int j = 0; j < 4; j++)
			{
				for (int k = 0; k < 4; k++)
				{
					p[j][k] = shPixelsY[tileCornerIdx + j * kTilePitch + k];
				}
			}

			shEdgeMap[edgeMapIdx] = vec4(GetEdgeMap(p, 0, 0));
			shEdgeMap[edgeMapIdx + 1] = vec4(GetEdgeMap(p, 0, 1));
			shEdgeMap[edgeMapIdx + kEdgeMapPitch] = vec4(GetEdgeMap(p, 1, 0));
			shEdgeMap[edgeMapIdx + kEdgeMapPitch + 1] = vec4(GetEdgeMap(p, 1, 1));
		}
	}

	LoadFilterBanksSh(int(threadIdx), NIS_THREAD_GROUP_SIZE);

	groupMemoryBarrier();
	barrier();

	// output coord within a tile
	const ivec2 pos = ivec2(uint(threadIdx) % uint(NIS_BLOCK_WIDTH), uint(threadIdx) / uint(NIS_BLOCK_WIDTH));
	// x coord inside the output image
	const int dstX = dstBlockX + pos.x;
	// x coord inside the input image
	const float srcX = (0.5f + dstX) * kScaleX - 0.5f;
	// nearest integer part
	const int px = int(floor(srcX) - srcBlockStartX);
	// fractional part
	const float fx = srcX - floor(srcX);
	// discretized phase
	const int fx_int = int(fx * kPhaseCount);

	for (int k = 0; k < NIS_BLOCK_WIDTH * NIS_BLOCK_HEIGHT / NIS_THREAD_GROUP_SIZE; ++k)
	{
		// y coord inside the output image
		const int dstY = dstBlockY + pos.y + k * (NIS_THREAD_GROUP_SIZE / NIS_BLOCK_WIDTH);
		// y coord inside the input image
		const float srcY = (0.5f + dstY) * kScaleY - 0.5f;

		// nearest integer part
		const int py = int(floor(srcY) - srcBlockStartY);
		// fractional part
		const float fy = srcY - floor(srcY);
		// discretized phase
		const int fy_int = int(fy * kPhaseCount);

		// generate weights for directional filters
		const int startEdgeMapIdx = py * kEdgeMapPitch + px;
		vec4 edge[2][2];
		for (int i = 0; i < 2; i++)
		{
			for (int 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 vec4 w = GetInterpEdgeMap(edge, fx, fy) * NIS_SCALE_INT;

		// load 6x6 support to regs
		const int startTileIdx = py * kTilePitch + px;
		float p[6][6];
		{
			for (int i = 0; i < 6; ++i)
			{
				for (int j = 0; j < 6; ++j)
				{
					p[i][j] = shPixelsY[startTileIdx + i * kTilePitch + j];
				}
			}
		}

		// weigth for luma
		const float baseWeight = NIS_SCALE_FLOAT - w.x - w.y - w.z - w.w;

		// final luma is a weighted product of directional & normal filters
		float 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);

		// do bilinear tap for luma upscaling
		vec4 op = HOOKED_tex(vec2((srcX + 0.5f) * kSrcNormX, (srcY + 0.5f) * kSrcNormY));

		const float corr = opY * (1.0f / NIS_SCALE_FLOAT) - op.r;
		op.x += corr;

		imageStore(out_image, ivec2(dstX, dstY), op);
	}
}

//!TEXTURE coef_scaler
//!SIZE 2 64
//!FORMAT rgba32f
//!FILTER NEAREST
00000000000000000000803f0000000000000000000000000000000000000000ed0d3e3ba91350bc0000803fd044583c89d25ebb0000000000000000000000003b70ce3b16fbcbbcb29d7f3f645ddc3c89d2debb000000000000000000000000e02d103c98dd13bda4df7e3fe7fb293d55c128bc6f12033a00000000000000005bb13f3c812642bd5b427e3ff931663d1ea768bc6f12033a00000000000000001ea7683cfaed6bbdfb5c7d3fbc05923d744694bc6f12033a0000000000000000b9fc873c03098abda3017c3fc5feb23d5839b4bc6f12833a0000000000000000075f983cbf0e9cbdfa7e7a3ff4fdd43dd044d8bca69bc43a00000000000000009eefa73c7b14aebdde02793f22fdf63d4850fcbc6f12033b0000000000000000ec51b83ced0dbebd22fd763f4d840d3ee02d10bd52491d3b0000000000000000efc9c33c5f07cebdb81e753f098a1f3e9c3322bded0d3e3b0000000000000000a913d03c88f4dbbd01de723fa1f8313e7dd033bd89d25e3b0000000000000000d044d83cf90fe9bd4e62703f9487453e82e247bde02d903b0000000000000000d3bce33cb3eaf3bd849e6d3f50fc583e88f45bbd2e90a03b0000000000000000faedeb3cdbf9febd83c06a3f448b6c3e44fa6dbdca54c13b00000000000000008e06f03c6f8104be82e2673f1283803e250681bd3b70ce3b0000000000000000b537f83ccc7f08be6f81643f83c08a3e280f8bbdd734ef3b00000000000000004850fc3c16fb0bbe97ff603f7d3f953e2b1895bdb9fc073c00000000000000004850fc3c60760fbe849e5d3f77be9f3ec0ec9ebd075f183c00000000000000006e34003dbc0512beecc0593ffa7eaa3ec3f5a8bd55c1283c00000000000000006e34003dcff713bec6dc553f7d3fb53ec5feb2bd3480373c00000000000000006e34003d068115bea5bd513f8941c03ec807bdbd82e2473c00000000000000006e34003d2b8716befb5c4d3f0c02cb3ea60ac6bdf775603c00000000000000004850fc3c197317be151d493fa245d63ea913d0bdd7346f3c0000000000000000b537f83c197317be34a2443f933ae13e8716d9bd24977f3c0000000000000000211ff43c197317bec520403f9f3cec3e6519e2bdb9fc873c00000000000000008e06f03c2b8716be083d3b3fab3ef73ed5e7eabd7446943c0000000000000000faedeb3c068115be143f363f2041013ffc18f3bde3a59b3c0000000000000000d3bce33ccff713bee561313f27c2063fb515fbbd0ad7a33c000000000000000040a4df3cce8812be68222c3f2d430c3f250601bec520b03c0000000000000000d044d83c857c10bee71d273fa5bd113f810405beec51b83c0000000000000000a913d03c5f070ebe6ade213fe71d173fb9fc07be5bb1bf3c000000000000000082e2c73cf1f40abe287e1c3f287e1c3ff1f40abe82e2c73c00000000000000005bb1bf3cb9fc07bee71d173f6ade213f5f070ebea913d03c0000000000000000ec51b83c810405bea5bd113fe71d273f857c10bed044d83c0000000000000000c520b03c250601be2d430c3f68222c3fce8812be40a4df3c00000000000000000ad7a33cb515fbbd27c2063fe561313fcff713bed3bce33c0000000000000000e3a59b3cfc18f3bd2041013f143f363f068115befaedeb3c00000000000000007446943cd5e7eabdab3ef73e083d3b3f2b8716be8e06f03c0000000000000000b9fc873c6519e2bd9f3cec3ec520403f197317be211ff43c000000000000000024977f3c8716d9bd933ae13e34a2443f197317beb537f83c0000000000000000d7346f3ca913d0bda245d63e151d493f197317be4850fc3c0000000000000000f775603ca60ac6bd0c02cb3efb5c4d3f2b8716be6e34003d000000000000000082e2473cc807bdbd8941c03ea5bd513f068115be6e34003d00000000000000003480373cc5feb2bd7d3fb53ec6dc553fcff713be6e34003d000000000000000055c1283cc3f5a8bdfa7eaa3eecc0593fbc0512be6e34003d0000000000000000075f183cc0ec9ebd77be9f3e849e5d3f60760fbe4850fc3c0000000000000000b9fc073c2b1895bd7d3f953e97ff603f16fb0bbe4850fc3c0000000000000000d734ef3b280f8bbd83c08a3e6f81643fcc7f08beb537f83c00000000000000003b70ce3b250681bd1283803e82e2673f6f8104be8e06f03c0000000000000000ca54c13b44fa6dbd448b6c3e83c06a3fdbf9febdfaedeb3c00000000000000002e90a03b88f45bbd50fc583e849e6d3fb3eaf3bdd3bce33c0000000000000000e02d903b82e247bd9487453e4e62703ff90fe9bdd044d83c000000000000000089d25e3b7dd033bda1f8313e01de723f88f4dbbda913d03c0000000000000000ed0d3e3b9c3322bd098a1f3eb81e753f5f07cebdefc9c33c000000000000000052491d3be02d10bd4d840d3e22fd763fed0dbebdec51b83c00000000000000006f12033b4850fcbc22fdf63dde02793f7b14aebd9eefa73c0000000000000000a69bc43ad044d8bcf4fdd43dfa7e7a3fbf0e9cbd075f983c00000000000000006f12833a5839b4bcc5feb23da3017c3f03098abdb9fc873c00000000000000006f12033a744694bcbc05923dfb5c7d3ffaed6bbd1ea7683c00000000000000006f12033a1ea768bcf931663d5b427e3f812642bd5bb13f3c00000000000000006f12033a55c128bce7fb293da4df7e3f98dd13bde02d103c00000000000000000000000089d2debb645ddc3cb29d7f3f16fbcbbc3b70ce3b00000000000000000000000089d25ebbd044583c0000803fa91350bced0d3e3b0000000000000000

//!TEXTURE coef_usm
//!SIZE 2 64
//!FORMAT rgba32f
//!FILTER NEAREST
0000000027a019bf27a0993f27a019bf00000000000000000000000000000000ed0d3e3b1ac01bbf006f993fe71d17bfed0d3ebb0000000000000000000000002e90a03bfb5c1dbff90f993fe63f14bf89d2debb6f12033a0000000000000000d734ef3b1b9e1ebf2731983fd3de10bf55c128bc000000000000000000000000075f183cb29d1fbfcb10973fff210dbffe6577bc0000000000000000000000003480373c4e6220bf48bf953fde0209bf77be9fbc6f12033a000000000000000082e2473c128320bfe63f943f34a204bf3d2cd4bc6f12033a00000000000000001ea7683cd3de20bfbe9f923fc52000bfdcd701bd6f12033a000000000000000024977f3c4e6220bfa54e903f7d3ff5be091b1ebd6f12033a0000000000000000b9fc873cb29d1fbf6ff08d3fe7fbe9be5af539bd6f12833a0000000000000000e02d903c20631ebf4f408b3fe483debe3ee859bd6f12833a00000000000000007446943cff211dbf696f883fbc05d2be6de77bbda69bc43a0000000000000000e3a59b3ccc5d1bbf1b2f853ff8c2c4be4df38ebda69bc43a000000000000000077be9f3cecc019bf910f823f22fdb6be5305a3bd6f12033b00000000000000000ad7a33cbec117bfc4427d3f423ea8be10e9b7bd52491d3b00000000000000000ad7a33cf4fd14bf7d3f753f50fc98be3b01cdbded0d3e3b00000000000000000ad7a33c05a312bf0de06d3f5eba89be6519e2bd89d25e3b00000000000000000ad7a33c3bdf0fbf8fc2653fb37b72beb515fbbd24977f3b00000000000000009eefa73cb1bf0cbfc4425d3faa8251bef08509bee02d903b00000000000000000ad7a33c637f09bf6f81543f69002fbe190416be2e90a03b000000000000000077be9f3c4f1e06bfcc5d4b3f16fb0bbe418222be7cf2b03b000000000000000077be9f3c3cbd02bf4182423fce19d1bda08930beca54c13b0000000000000000e3a59b3c5b42febe151d393f4bea84bddbf93ebe3b70ce3b0000000000000000075f983c99bbf6bef2412f3ffaedebbc287e4cbe89d2de3b0000000000000000075f983c6900efbe4260253f075f183cac8b5bbed734ef3b0000000000000000e02d903c27c2e6be0c021b3fca32443dfa7e6abeb9fc073c00000000000000004d158c3c77bedfbea5bd113f57ecaf3d6c787abee02d103c000000000000000026e4833c22fdd6beab3e073f1283003e810485be2e90203c000000000000000026e4833cf241cfbe7502fa3ed578293e3b018dbe55c1283c0000000000000000fe65773cb003c7be143fe63e97ff503ef4fd94be0e4f2f3c00000000000000001ea7683cd200bebe857cd03e6c787a3eadfa9cbe5bb13f3c0000000000000000f775603c1904b6bea301bc3ebc05923e0b46a5be82e2473c0000000000000000d044583cd6c5adbeb003a73eb003a73ed6c5adbed044583c000000000000000082e2473c0b46a5bebc05923ea301bc3e1904b6bef775603c00000000000000005bb13f3cadfa9cbe6c787a3e857cd03ed200bebe1ea7683c00000000000000000e4f2f3cf4fd94be97ff503e143fe63eb003c7befe65773c000000000000000055c1283c3b018dbed578293e7502fa3ef241cfbe26e4833c00000000000000002e90203c810485be1283003eab3e073f22fdd6be26e4833c0000000000000000e02d103c6c787abe57ecaf3da5bd113f77bedfbe4d158c3c0000000000000000b9fc073cfa7e6abeca32443d0c021b3f27c2e6bee02d903c0000000000000000d734ef3bac8b5bbe075f183c4260253f6900efbe075f983c000000000000000089d2de3b287e4cbefaedebbcf2412f3f99bbf6be075f983c00000000000000003b70ce3bdbf93ebe4bea84bd151d393f5b42febee3a59b3c0000000000000000ca54c13ba08930bece19d1bd4182423f3cbd02bf77be9f3c00000000000000007cf2b03b418222be16fb0bbecc5d4b3f4f1e06bf77be9f3c00000000000000002e90a03b190416be69002fbe6f81543f637f09bf0ad7a33c0000000000000000e02d903bf08509beaa8251bec4425d3fb1bf0cbf9eefa73c000000000000000024977f3bb515fbbdb37b72be8fc2653f3bdf0fbf0ad7a33c000000000000000089d25e3b6519e2bd5eba89be0de06d3f05a312bf0ad7a33c0000000000000000ed0d3e3b3b01cdbd50fc98be7d3f753ff4fd14bf0ad7a33c000000000000000052491d3b10e9b7bd423ea8bec4427d3fbec117bf0ad7a33c00000000000000006f12033b5305a3bd22fdb6be910f823fecc019bf77be9f3c0000000000000000a69bc43a4df38ebdf8c2c4be1b2f853fcc5d1bbfe3a59b3c0000000000000000a69bc43a6de77bbdbc05d2be696f883fff211dbf7446943c00000000000000006f12833a3ee859bde483debe4f408b3f20631ebfe02d903c00000000000000006f12833a5af539bde7fbe9be6ff08d3fb29d1fbfb9fc873c00000000000000006f12033a091b1ebd7d3ff5bea54e903f4e6220bf24977f3c00000000000000006f12033adcd701bdc52000bfbe9f923fd3de20bf1ea7683c00000000000000006f12033a3d2cd4bc34a204bfe63f943f128320bf82e2473c00000000000000006f12033a77be9fbcde0209bf48bf953f4e6220bf3480373c000000000000000000000000fe6577bcff210dbfcb10973fb29d1fbf075f183c00000000000000000000000055c128bcd3de10bf2731983f1b9e1ebfd734ef3b00000000000000006f12033a89d2debbe63f14bff90f993ffb5c1dbf2e90a03b000000000000000000000000ed0d3ebbe71d17bf006f993f1ac01bbfed0d3e3b0000000000000000

NVSharpen.glsl

// 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 v1.0.2 by NVIDIA
// ported to mpv by agyild

// Changelog
// Made it directly operate on LUMA plane, since the original shader was operating
// on LUMA by deriving it from RGB.

//!HOOK LUMA
//!BIND HOOKED
//!DESC NVIDIA Image Sharpening v1.0.2
//!COMPUTE 32 32 256 1
//!WHEN OUTPUT.w OUTPUT.h * LUMA.w LUMA.h * / 1.0 > ! OUTPUT.w OUTPUT.h * LUMA.w LUMA.h * / 1.0 < ! *

// User variables
#define SHARPNESS 0.25 // Amount of sharpening. 0.0 to 1.0.
#define NIS_THREAD_GROUP_SIZE 256 // May be set to 128 for better performance on NVIDIA hardware, otherwise set to 256. Don't forget to modify the COMPUTE directive accordingly as well (e.g., COMPUTE 32 32 128 1).
#define NIS_HDR_MODE 0 // Must be set to 1 for content with PQ colorspace. 0 or 1.

// Constant variables
#define NIS_BLOCK_WIDTH 32
#define NIS_BLOCK_HEIGHT 32
#define kSupportSize 5
#define kNumPixelsX (NIS_BLOCK_WIDTH + kSupportSize + 1)
#define kNumPixelsY (NIS_BLOCK_HEIGHT + kSupportSize + 1)
const float sharpen_slider = clamp(SHARPNESS, 0.0f, 1.0f) - 0.5f;
const float MaxScale = (sharpen_slider >= 0.0f) ? 1.25f : 1.75f;
const float MinScale = (sharpen_slider >= 0.0f) ? 1.25f : 1.0f;
const float LimitScale = (sharpen_slider >= 0.0f) ? 1.25f : 1.0f;
const float kDetectRatio = 2 * 1127.f / 1024.f;
const float kDetectThres = (bool(NIS_HDR_MODE) ? 32.0f : 64.0f) / 1024.0f;
const float kMinContrastRatio = bool(NIS_HDR_MODE) ? 1.5f : 2.0f;
const float kMaxContrastRatio = bool(NIS_HDR_MODE) ? 5.0f : 10.0f;
const float kSharpStartY = bool(NIS_HDR_MODE) ? 0.35f : 0.45f;
const float kSharpEndY = bool(NIS_HDR_MODE) ? 0.55f : 0.9f;
const float kSharpStrengthMin = max(0.0f, 0.4f + sharpen_slider * MinScale * (bool(NIS_HDR_MODE) ? 1.1f : 1.2));
const float kSharpStrengthMax = ((bool(NIS_HDR_MODE) ? 2.2f : 1.6f) + sharpen_slider * MaxScale * 1.8f);
const float kSharpLimitMin = max((bool(NIS_HDR_MODE) ? 0.06f :0.1f), (bool(NIS_HDR_MODE) ? 0.1f : 0.14f) + sharpen_slider * LimitScale * (bool(NIS_HDR_MODE) ? 0.28f : 0.32f)); //
const float kSharpLimitMax = ((bool(NIS_HDR_MODE) ? 0.6f : 0.5f) + sharpen_slider * LimitScale * 0.6f);
const float kRatioNorm = 1.0f / (kMaxContrastRatio - kMinContrastRatio);
const float kSharpScaleY = 1.0f / (kSharpEndY - kSharpStartY);
const float kSharpStrengthScale = kSharpStrengthMax - kSharpStrengthMin;
const float kSharpLimitScale = kSharpLimitMax - kSharpLimitMin;
const float kContrastBoost = 1.0f;
const float kEps = 1.0f / 255.0f;
#define kSrcNormX HOOKED_pt.x
#define kSrcNormY HOOKED_pt.y
#define kDstNormX kSrcNormX
#define kDstNormY kSrcNormY

// HLSL to GLSL macros
#define saturate(x) clamp(x, 0, 1)
#define lerp(a, b, x) mix(a, b, x)

// CS Shared variables
shared float shPixelsY[kNumPixelsY][kNumPixelsX];

// Shader code

vec4 GetEdgeMap(float p[5][5], int i, int j) {
	const float 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 float 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 float 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 float 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 float g_0_90_max = max(g_0, g_90);
	const float g_0_90_min = min(g_0, g_90);
	const float g_45_135_max = max(g_45, g_135);
	const float g_45_135_min = min(g_45, g_135);

	float e_0_90 = 0;
	float e_45_135 = 0;

    if (g_0_90_max + g_45_135_max == 0)
    {
        return vec4(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;

    bool 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);
    bool 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);
    bool c_g_0_90 = g_0_90_max == g_0;
    bool c_g_45_135 = g_45_135_max == g_45;

    float f_e_0_90 = (c_0_90 && c_45_135) ? e_0_90 : 1.0f;
    float f_e_45_135 = (c_0_90 && c_45_135) ? e_45_135 : 1.0f;

    float weight_0 = (c_0_90 && c_g_0_90) ? f_e_0_90 : 0.0f;
    float weight_90 = (c_0_90 && !c_g_0_90) ? f_e_0_90 : 0.0f;
    float weight_45 = (c_45_135 && c_g_45_135) ? f_e_45_135 : 0.0f;
    float weight_135 = (c_45_135 && !c_g_45_135) ? f_e_45_135 : 0.0f;

	return vec4(weight_0, weight_90, weight_45, weight_135);
}

float CalcLTIFast(const float y[5]) {
	const float a_min = min(min(y[0], y[1]), y[2]);
	const float a_max = max(max(y[0], y[1]), y[2]);

	const float b_min = min(min(y[2], y[3]), y[4]);
	const float b_max = max(max(y[2], y[3]), y[4]);

	const float a_cont = a_max - a_min;
	const float b_cont = b_max - b_min;

	const float cont_ratio = max(a_cont, b_cont) / (min(a_cont, b_cont) + kEps);
	return (1.0f - saturate((cont_ratio - kMinContrastRatio) * kRatioNorm)) * kContrastBoost;
}

float EvalUSM(const float pxl[5], const float sharpnessStrength, const float sharpnessLimit) {
	// USM profile
	float 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;
}

vec4 GetDirUSM(const float p[5][5]) {
	// sharpness boost & limit are the same for all directions
	const float scaleY = 1.0f - saturate((p[2][2] - kSharpStartY) * kSharpScaleY);
	// scale the ramp to sharpen as a function of luma
	const float sharpnessStrength = scaleY * kSharpStrengthScale + kSharpStrengthMin;
	// scale the ramp to limit USM as a function of luma
	const float sharpnessLimit = (scaleY * kSharpLimitScale + kSharpLimitMin) * p[2][2];

	vec4 rval;
	// 0 deg filter
	float interp0Deg[5];
	{
		for (int i = 0; i < 5; ++i)
		{
			interp0Deg[i] = p[i][2];
		}
	}

	rval.x = EvalUSM(interp0Deg, sharpnessStrength, sharpnessLimit);

	// 90 deg filter
	float interp90Deg[5];
	{
		for (int i = 0; i < 5; ++i)
		{
			interp90Deg[i] = p[2][i];
		}
	}

	rval.y = EvalUSM(interp90Deg, sharpnessStrength, sharpnessLimit);

	//45 deg filter
	float 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
	float 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;
}

void hook() {
	uvec2 blockIdx = gl_WorkGroupID.xy;
	uint threadIdx = gl_LocalInvocationID.x;

	const int dstBlockX = int(NIS_BLOCK_WIDTH * blockIdx.x);
	const int dstBlockY = int(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 float kShift = 0.5f - kSupportSize / 2;

	for (int i = int(threadIdx) * 2; i < kNumPixelsX * kNumPixelsY / 2; i += NIS_THREAD_GROUP_SIZE * 2) {
		uvec2 pos = uvec2(uint(i) % uint(kNumPixelsX), uint(i) / uint(kNumPixelsX) * 2);

		for (int dy = 0; dy < 2; dy++) {
			for (int dx = 0; dx < 2; dx++) {
				const float tx = (dstBlockX + pos.x + dx + kShift) * kSrcNormX;
				const float ty = (dstBlockY + pos.y + dy + kShift) * kSrcNormY;
				const float px = HOOKED_tex(vec2(tx, ty)).r;
				shPixelsY[pos.y + dy][pos.x + dx] = px;
			}
		}
	}

	groupMemoryBarrier();
	barrier();

	for (int k = int(threadIdx); k < NIS_BLOCK_WIDTH * NIS_BLOCK_HEIGHT; k += NIS_THREAD_GROUP_SIZE)
	{
		const ivec2 pos = ivec2(uint(k) % uint(NIS_BLOCK_WIDTH), uint(k) / uint(NIS_BLOCK_WIDTH));

		// load 5x5 support to regs
		float p[5][5];

		for (int i = 0; i < 5; ++i)
		{
			for (int j = 0; j < 5; ++j)
			{
				p[i][j] = shPixelsY[pos.y + i][pos.x + j];
			}
		}

		// get directional filter bank output
		vec4 dirUSM = GetDirUSM(p);

		// generate weights for directional filters
		vec4 w = GetEdgeMap(p, kSupportSize / 2 - 1, kSupportSize / 2 - 1);

		// final USM is a weighted sum filter outputs
		const float usmY = (dirUSM.x * w.x + dirUSM.y * w.y + dirUSM.z * w.z + dirUSM.w * w.w);

		// do bilinear tap and correct luma texel so it produces new sharpened luma
		const int dstX = dstBlockX + pos.x;
		const int dstY = dstBlockY + pos.y;

		vec4 op = HOOKED_tex(vec2((dstX + 0.5f) * kDstNormX, (dstY + 0.5f) * kDstNormY));
		op.x += usmY;

		imageStore(out_image, ivec2(dstX, dstY), op);
	}
}

You can disable shaders on rotate like this:

[profile-name]
profile-cond=(p["video-params/rotate"] == 0)
glsl-shaders-clr
glsl-shader="~~/shaders/NVScaler.glsl"

Use the glsl-shaders-clr part wisely ;)

How do i reduce the ringing artifacts? FSR has built in de-ringing.

How do i reduce the ringing artifacts? FSR has built in de-ringing.

NIS does not have built-in deringing. You might want to try adding antiring.hook after NVScaler.glsl and configuring the radius parameter for both passes to achieve the desired effect.