avfilter/vf_yadif_cuda: CUDA accelerated yadif deinterlacer

This is a cuda implementation of yadif, which gives us a way to
do deinterlacing when using the nvdec hwaccel. In that scenario
we don't have access to the nvidia deinterlacer.

1 files changed, 296 insertions, 0 deletions

diff --git a/libavfilter/vf_yadif_cuda.cu b/libavfilter/vf_yadif_cuda.cu
new file mode 100644
index 0000000..65a902c
--- /dev/null
+++ b/libavfilter/vf_yadif_cuda.cu
@@ -0,0 +1,296 @@
+/*
+ * Copyright (C) 2018 Philip Langdale <philipl@overt.org>
+ *
+ * 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
+ */
+
+template<typename T>
+__inline__ __device__ 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)
+{
+    int spatial_pred = (d + k)/2;
+    int spatial_score = abs(c - j) + abs(d - k) + abs(e - l);
+
+    int 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;
+}
+
+__inline__ __device__ int max3(int a, int b, int c)
+{
+    int x = max(a, b);
+    return max(x, c);
+}
+
+__inline__ __device__ int min3(int a, int b, int c)
+{
+    int x = min(a, b);
+    return min(x, c);
+}
+
+template<typename T>
+__inline__ __device__ 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)
+{
+    int p0 = (C + H) / 2;
+    int p1 = F;
+    int p2 = (D + I) / 2;
+    int p3 = G;
+    int p4 = (E + J) / 2;
+
+    int tdiff0 = abs(D - I);
+    int tdiff1 = (abs(A - F) + abs(B - G)) / 2;
+    int tdiff2 = (abs(K - F) + abs(G - L)) / 2;
+
+    int diff = max3(tdiff0, tdiff1, tdiff2);
+
+    if (!skip_check) {
+      int maxi = max3(p2 - p3, p2 - p1, min(p0 - p1, p4 - p3));
+      int mini = min3(p2 - p3, p2 - p1, max(p0 - p1, p4 - p3));
+      diff = max3(diff, mini, -maxi);
+    }
+
+    if (spatial_pred > p2 + diff) {
+      spatial_pred = p2 + diff;
+    }
+    if (spatial_pred < p2 - diff) {
+      spatial_pred = p2 - diff;
+    }
+
+    return spatial_pred;
+}
+
+template<typename T>
+__inline__ __device__ void yadif_single(T *dst,
+                                        cudaTextureObject_t prev,
+                                        cudaTextureObject_t cur,
+                                        cudaTextureObject_t next,
+                                        int dst_width, int dst_height, int dst_pitch,
+                                        int src_width, int src_height,
+                                        int parity, int tff, bool skip_spatial_check)
+{
+    // Identify location
+    int xo = blockIdx.x * blockDim.x + threadIdx.x;
+    int yo = blockIdx.y * blockDim.y + threadIdx.y;
+
+    if (xo >= dst_width || yo >= dst_height) {
+        return;
+    }
+
+    // Don't modify the primary field
+    if (yo % 2 == parity) {
+      dst[yo*dst_pitch+xo] = tex2D<T>(cur, xo, yo);
+      return;
+    }
+
+    // Calculate spatial prediction
+    T a = tex2D<T>(cur, xo - 3, yo - 1);
+    T b = tex2D<T>(cur, xo - 2, yo - 1);
+    T c = tex2D<T>(cur, xo - 1, yo - 1);
+    T d = tex2D<T>(cur, xo - 0, yo - 1);
+    T e = tex2D<T>(cur, xo + 1, yo - 1);
+    T f = tex2D<T>(cur, xo + 2, yo - 1);
+    T g = tex2D<T>(cur, xo + 3, yo - 1);
+
+    T h = tex2D<T>(cur, xo - 3, yo + 1);
+    T i = tex2D<T>(cur, xo - 2, yo + 1);
+    T j = tex2D<T>(cur, xo - 1, yo + 1);
+    T k = tex2D<T>(cur, xo - 0, yo + 1);
+    T l = tex2D<T>(cur, xo + 1, yo + 1);
+    T m = tex2D<T>(cur, xo + 2, yo + 1);
+    T n = tex2D<T>(cur, xo + 3, yo + 1);
+
+    T spatial_pred =
+        spatial_predictor(a, b, c, d, e, f, g, h, i, j, k, l, m, n);
+
+    // Calculate temporal prediction
+    int is_second_field = !(parity ^ tff);
+
+    cudaTextureObject_t prev2 = prev;
+    cudaTextureObject_t prev1 = is_second_field ? cur : prev;
+    cudaTextureObject_t next1 = is_second_field ? next : cur;
+    cudaTextureObject_t next2 = next;
+
+    T A = tex2D<T>(prev2, xo,  yo - 1);
+    T B = tex2D<T>(prev2, xo,  yo + 1);
+    T C = tex2D<T>(prev1, xo,  yo - 2);
+    T D = tex2D<T>(prev1, xo,  yo + 0);
+    T E = tex2D<T>(prev1, xo,  yo + 2);
+    T F = tex2D<T>(cur,   xo,  yo - 1);
+    T G = tex2D<T>(cur,   xo,  yo + 1);
+    T H = tex2D<T>(next1, xo,  yo - 2);
+    T I = tex2D<T>(next1, xo,  yo + 0);
+    T J = tex2D<T>(next1, xo,  yo + 2);
+    T K = tex2D<T>(next2, xo,  yo - 1);
+    T L = tex2D<T>(next2, xo,  yo + 1);
+
+    spatial_pred = temporal_predictor(A, B, C, D, E, F, G, H, I, J, K, L,
+                                      spatial_pred, skip_spatial_check);
+
+    dst[yo*dst_pitch+xo] = spatial_pred;
+}
+
+template <typename T>
+__inline__ __device__ void yadif_double(T *dst,
+                                        cudaTextureObject_t prev,
+                                        cudaTextureObject_t cur,
+                                        cudaTextureObject_t next,
+                                        int dst_width, int dst_height, int dst_pitch,
+                                        int src_width, int src_height,
+                                        int parity, int tff, bool skip_spatial_check)
+{
+    int xo = blockIdx.x * blockDim.x + threadIdx.x;
+    int yo = blockIdx.y * blockDim.y + threadIdx.y;
+
+    if (xo >= dst_width || yo >= dst_height) {
+        return;
+    }
+
+    if (yo % 2 == parity) {
+      // Don't modify the primary field
+      dst[yo*dst_pitch+xo] = tex2D<T>(cur, xo, yo);
+      return;
+    }
+
+    T a = tex2D<T>(cur, xo - 3, yo - 1);
+    T b = tex2D<T>(cur, xo - 2, yo - 1);
+    T c = tex2D<T>(cur, xo - 1, yo - 1);
+    T d = tex2D<T>(cur, xo - 0, yo - 1);
+    T e = tex2D<T>(cur, xo + 1, yo - 1);
+    T f = tex2D<T>(cur, xo + 2, yo - 1);
+    T g = tex2D<T>(cur, xo + 3, yo - 1);
+
+    T h = tex2D<T>(cur, xo - 3, yo + 1);
+    T i = tex2D<T>(cur, xo - 2, yo + 1);
+    T j = tex2D<T>(cur, xo - 1, yo + 1);
+    T k = tex2D<T>(cur, xo - 0, yo + 1);
+    T l = tex2D<T>(cur, xo + 1, yo + 1);
+    T m = tex2D<T>(cur, xo + 2, yo + 1);
+    T n = tex2D<T>(cur, xo + 3, yo + 1);
+
+    T spatial_pred = {
+        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) };
+
+    // Calculate temporal prediction
+    int is_second_field = !(parity ^ tff);
+
+    cudaTextureObject_t prev2 = prev;
+    cudaTextureObject_t prev1 = is_second_field ? cur : prev;
+    cudaTextureObject_t next1 = is_second_field ? next : cur;
+    cudaTextureObject_t next2 = next;
+
+    T A = tex2D<T>(prev2, xo,  yo - 1);
+    T B = tex2D<T>(prev2, xo,  yo + 1);
+    T C = tex2D<T>(prev1, xo,  yo - 2);
+    T D = tex2D<T>(prev1, xo,  yo + 0);
+    T E = tex2D<T>(prev1, xo,  yo + 2);
+    T F = tex2D<T>(cur,   xo,  yo - 1);
+    T G = tex2D<T>(cur,   xo,  yo + 1);
+    T H = tex2D<T>(next1, xo,  yo - 2);
+    T I = tex2D<T>(next1, xo,  yo + 0);
+    T J = tex2D<T>(next1, xo,  yo + 2);
+    T K = tex2D<T>(next2, xo,  yo - 1);
+    T L = tex2D<T>(next2, xo,  yo + 1);
+
+    spatial_pred = {
+        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_spatial_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_spatial_check) };
+
+    dst[yo*dst_pitch+xo] = spatial_pred;
+}
+
+extern "C" {
+
+__global__ void yadif_uchar(unsigned char *dst,
+                            cudaTextureObject_t prev,
+                            cudaTextureObject_t cur,
+                            cudaTextureObject_t next,
+                            int dst_width, int dst_height, int dst_pitch,
+                            int src_width, int src_height,
+                            int parity, int tff, bool skip_spatial_check)
+{
+    yadif_single(dst, prev, cur, next,
+                 dst_width, dst_height, dst_pitch,
+                 src_width, src_height,
+                 parity, tff, skip_spatial_check);
+}
+
+__global__ void yadif_ushort(unsigned short *dst,
+                            cudaTextureObject_t prev,
+                            cudaTextureObject_t cur,
+                            cudaTextureObject_t next,
+                            int dst_width, int dst_height, int dst_pitch,
+                            int src_width, int src_height,
+                            int parity, int tff, bool skip_spatial_check)
+{
+    yadif_single(dst, prev, cur, next,
+                 dst_width, dst_height, dst_pitch,
+                 src_width, src_height,
+                 parity, tff, skip_spatial_check);
+}
+
+__global__ void yadif_uchar2(uchar2 *dst,
+                            cudaTextureObject_t prev,
+                            cudaTextureObject_t cur,
+                            cudaTextureObject_t next,
+                            int dst_width, int dst_height, int dst_pitch,
+                            int src_width, int src_height,
+                            int parity, int tff, bool skip_spatial_check)
+{
+    yadif_double(dst, prev, cur, next,
+                 dst_width, dst_height, dst_pitch,
+                 src_width, src_height,
+                 parity, tff, skip_spatial_check);
+}
+
+__global__ void yadif_ushort2(ushort2 *dst,
+                            cudaTextureObject_t prev,
+                            cudaTextureObject_t cur,
+                            cudaTextureObject_t next,
+                            int dst_width, int dst_height, int dst_pitch,
+                            int src_width, int src_height,
+                            int parity, int tff, bool skip_spatial_check)
+{
+    yadif_double(dst, prev, cur, next,
+                 dst_width, dst_height, dst_pitch,
+                 src_width, src_height,
+                 parity, tff, skip_spatial_check);
+}
+
+} /* extern "C" */