/* * Copyright (C) 2018 Philip Langdale * 2020 Aman Karmani * 2020 Stefan Dyulgerov * * This file is part of FFmpeg. * * FFmpeg is free software; you can redistribute it and/or * modify it under the terms of the GNU Lesser General Public * License as published by the Free Software Foundation; either * version 2.1 of the License, or (at your option) any later version. * * FFmpeg is distributed in the hope that it will be useful, * but WITHOUT ANY WARRANTY; without even the implied warranty of * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU * Lesser General Public License for more details. * * You should have received a copy of the GNU Lesser General Public * License along with FFmpeg; if not, write to the Free Software * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA */ #include #include #include using namespace metal; /* * Version compat shims */ #if __METAL_VERSION__ < 210 #define max3(x, y, z) max(x, max(y, z)) #define min3(x, y, z) min(x, min(y, z)) #endif /* * Parameters */ struct deintParams { uint channels; uint parity; uint tff; bool is_second_field; bool skip_spatial_check; int field_mode; }; /* * Texture access helpers */ #define accesstype access::sample constexpr sampler s(coord::pixel); template T tex2D(texture2d tex, uint x, uint y) { return tex.sample(s, float2(x, y)).x; } template <> float2 tex2D(texture2d tex, uint x, uint y) { return tex.sample(s, float2(x, y)).xy; } template T tex2D(texture2d tex, uint x, uint y) { return tex.read(uint2(x, y)).x; } template <> float2 tex2D(texture2d tex, uint x, uint y) { return tex.read(uint2(x, y)).xy; } /* * YADIF helpers */ template T spatial_predictor(T a, T b, T c, T d, T e, T f, T g, T h, T i, T j, T k, T l, T m, T n) { T spatial_pred = (d + k)/2; T spatial_score = abs(c - j) + abs(d - k) + abs(e - l); T score = abs(b - k) + abs(c - l) + abs(d - m); if (score < spatial_score) { spatial_pred = (c + l)/2; spatial_score = score; score = abs(a - l) + abs(b - m) + abs(c - n); if (score < spatial_score) { spatial_pred = (b + m)/2; spatial_score = score; } } score = abs(d - i) + abs(e - j) + abs(f - k); if (score < spatial_score) { spatial_pred = (e + j)/2; spatial_score = score; score = abs(e - h) + abs(f - i) + abs(g - j); if (score < spatial_score) { spatial_pred = (f + i)/2; spatial_score = score; } } return spatial_pred; } template T temporal_predictor(T A, T B, T C, T D, T E, T F, T G, T H, T I, T J, T K, T L, T spatial_pred, bool skip_check) { T p0 = (C + H) / 2; T p1 = F; T p2 = (D + I) / 2; T p3 = G; T p4 = (E + J) / 2; T tdiff0 = abs(D - I); T tdiff1 = (abs(A - F) + abs(B - G)) / 2; T tdiff2 = (abs(K - F) + abs(G - L)) / 2; T diff = max3(tdiff0, tdiff1, tdiff2); if (!skip_check) { T maxi = max3(p2 - p3, p2 - p1, min(p0 - p1, p4 - p3)); T mini = min3(p2 - p3, p2 - p1, max(p0 - p1, p4 - p3)); diff = max3(diff, mini, -maxi); } return clamp(spatial_pred, p2 - diff, p2 + diff); } #define T float2 template <> T spatial_predictor(T a, T b, T c, T d, T e, T f, T g, T h, T i, T j, T k, T l, T m, T n) { return T( spatial_predictor(a.x, b.x, c.x, d.x, e.x, f.x, g.x, h.x, i.x, j.x, k.x, l.x, m.x, n.x), spatial_predictor(a.y, b.y, c.y, d.y, e.y, f.y, g.y, h.y, i.y, j.y, k.y, l.y, m.y, n.y) ); } template <> T temporal_predictor(T A, T B, T C, T D, T E, T F, T G, T H, T I, T J, T K, T L, T spatial_pred, bool skip_check) { return T( temporal_predictor(A.x, B.x, C.x, D.x, E.x, F.x, G.x, H.x, I.x, J.x, K.x, L.x, spatial_pred.x, skip_check), temporal_predictor(A.y, B.y, C.y, D.y, E.y, F.y, G.y, H.y, I.y, J.y, K.y, L.y, spatial_pred.y, skip_check) ); } #undef T /* * YADIF compute */ template T yadif_compute_spatial( texture2d cur, uint2 pos) { // Calculate spatial prediction T a = tex2D(cur, pos.x - 3, pos.y - 1); T b = tex2D(cur, pos.x - 2, pos.y - 1); T c = tex2D(cur, pos.x - 1, pos.y - 1); T d = tex2D(cur, pos.x - 0, pos.y - 1); T e = tex2D(cur, pos.x + 1, pos.y - 1); T f = tex2D(cur, pos.x + 2, pos.y - 1); T g = tex2D(cur, pos.x + 3, pos.y - 1); T h = tex2D(cur, pos.x - 3, pos.y + 1); T i = tex2D(cur, pos.x - 2, pos.y + 1); T j = tex2D(cur, pos.x - 1, pos.y + 1); T k = tex2D(cur, pos.x - 0, pos.y + 1); T l = tex2D(cur, pos.x + 1, pos.y + 1); T m = tex2D(cur, pos.x + 2, pos.y + 1); T n = tex2D(cur, pos.x + 3, pos.y + 1); return spatial_predictor(a, b, c, d, e, f, g, h, i, j, k, l, m, n); } template T yadif_compute_temporal( texture2d cur, texture2d prev2, texture2d prev1, texture2d next1, texture2d next2, T spatial_pred, bool skip_spatial_check, uint2 pos) { // Calculate temporal prediction T A = tex2D(prev2, pos.x, pos.y - 1); T B = tex2D(prev2, pos.x, pos.y + 1); T C = tex2D(prev1, pos.x, pos.y - 2); T D = tex2D(prev1, pos.x, pos.y + 0); T E = tex2D(prev1, pos.x, pos.y + 2); T F = tex2D(cur, pos.x, pos.y - 1); T G = tex2D(cur, pos.x, pos.y + 1); T H = tex2D(next1, pos.x, pos.y - 2); T I = tex2D(next1, pos.x, pos.y + 0); T J = tex2D(next1, pos.x, pos.y + 2); T K = tex2D(next2, pos.x, pos.y - 1); T L = tex2D(next2, pos.x, pos.y + 1); return temporal_predictor(A, B, C, D, E, F, G, H, I, J, K, L, spatial_pred, skip_spatial_check); } template T yadif( texture2d dst, texture2d prev, texture2d cur, texture2d next, constant deintParams& params, uint2 pos) { T spatial_pred = yadif_compute_spatial(cur, pos); if (params.is_second_field) { return yadif_compute_temporal(cur, prev, cur, next, next, spatial_pred, params.skip_spatial_check, pos); } else { return yadif_compute_temporal(cur, prev, prev, cur, next, spatial_pred, params.skip_spatial_check, pos); } } /* * Kernel dispatch */ kernel void deint( texture2d dst [[texture(0)]], texture2d prev [[texture(1)]], texture2d cur [[texture(2)]], texture2d next [[texture(3)]], constant deintParams& params [[buffer(4)]], uint2 pos [[thread_position_in_grid]]) { if ((pos.x >= dst.get_width()) || (pos.y >= dst.get_height())) { return; } // Don't modify the primary field if (pos.y % 2 == params.parity) { float4 in = cur.read(pos); dst.write(in, pos); return; } float2 pred; if (params.channels == 1) pred = float2(yadif(dst, prev, cur, next, params, pos)); else pred = yadif(dst, prev, cur, next, params, pos); dst.write(pred.xyyy, pos); }