git.sesse.net Git - ffmpeg/blob - libavfilter/vf_yadif_cuda.cu

   1 /*
   2  * Copyright (C) 2018 Philip Langdale <philipl@overt.org>
   3  *
   4  * This file is part of FFmpeg.
   5  *
   6  * FFmpeg is free software; you can redistribute it and/or
   7  * modify it under the terms of the GNU Lesser General Public
   8  * License as published by the Free Software Foundation; either
   9  * version 2.1 of the License, or (at your option) any later version.
  10  *
  11  * FFmpeg is distributed in the hope that it will be useful,
  12  * but WITHOUT ANY WARRANTY; without even the implied warranty of
  13  * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
  14  * Lesser General Public License for more details.
  15  *
  16  * You should have received a copy of the GNU Lesser General Public
  17  * License along with FFmpeg; if not, write to the Free Software
  18  * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
  19  */
  20
  21 template<typename T>
  22 __inline__ __device__ T spatial_predictor(T a, T b, T c, T d, T e, T f, T g,
  23                                           T h, T i, T j, T k, T l, T m, T n)
  24 {
  25     int spatial_pred = (d + k)/2;
  26     int spatial_score = abs(c - j) + abs(d - k) + abs(e - l);
  27
  28     int score = abs(b - k) + abs(c - l) + abs(d - m);
  29     if (score < spatial_score) {
  30         spatial_pred = (c + l)/2;
  31         spatial_score = score;
  32         score = abs(a - l) + abs(b - m) + abs(c - n);
  33         if (score < spatial_score) {
  34           spatial_pred = (b + m)/2;
  35           spatial_score = score;
  36         }
  37     }
  38     score = abs(d - i) + abs(e - j) + abs(f - k);
  39     if (score < spatial_score) {
  40         spatial_pred = (e + j)/2;
  41         spatial_score = score;
  42         score = abs(e - h) + abs(f - i) + abs(g - j);
  43         if (score < spatial_score) {
  44           spatial_pred = (f + i)/2;
  45           spatial_score = score;
  46         }
  47     }
  48     return spatial_pred;
  49 }
  50
  51 __inline__ __device__ int max3(int a, int b, int c)
  52 {
  53     int x = max(a, b);
  54     return max(x, c);
  55 }
  56
  57 __inline__ __device__ int min3(int a, int b, int c)
  58 {
  59     int x = min(a, b);
  60     return min(x, c);
  61 }
  62
  63 template<typename T>
  64 __inline__ __device__ T temporal_predictor(T A, T B, T C, T D, T E, T F,
  65                                            T G, T H, T I, T J, T K, T L,
  66                                            T spatial_pred, bool skip_check)
  67 {
  68     int p0 = (C + H) / 2;
  69     int p1 = F;
  70     int p2 = (D + I) / 2;
  71     int p3 = G;
  72     int p4 = (E + J) / 2;
  73
  74     int tdiff0 = abs(D - I);
  75     int tdiff1 = (abs(A - F) + abs(B - G)) / 2;
  76     int tdiff2 = (abs(K - F) + abs(G - L)) / 2;
  77
  78     int diff = max3(tdiff0, tdiff1, tdiff2);
  79
  80     if (!skip_check) {
  81       int maxi = max3(p2 - p3, p2 - p1, min(p0 - p1, p4 - p3));
  82       int mini = min3(p2 - p3, p2 - p1, max(p0 - p1, p4 - p3));
  83       diff = max3(diff, mini, -maxi);
  84     }
  85
  86     if (spatial_pred > p2 + diff) {
  87       spatial_pred = p2 + diff;
  88     }
  89     if (spatial_pred < p2 - diff) {
  90       spatial_pred = p2 - diff;
  91     }
  92
  93     return spatial_pred;
  94 }
  95
  96 template<typename T>
  97 __inline__ __device__ void yadif_single(T *dst,
  98                                         cudaTextureObject_t prev,
  99                                         cudaTextureObject_t cur,
 100                                         cudaTextureObject_t next,
 101                                         int dst_width, int dst_height, int dst_pitch,
 102                                         int src_width, int src_height,
 103                                         int parity, int tff, bool skip_spatial_check)
 104 {
 105     // Identify location
 106     int xo = blockIdx.x * blockDim.x + threadIdx.x;
 107     int yo = blockIdx.y * blockDim.y + threadIdx.y;
 108
 109     if (xo >= dst_width || yo >= dst_height) {
 110         return;
 111     }
 112
 113     // Don't modify the primary field
 114     if (yo % 2 == parity) {
 115       dst[yo*dst_pitch+xo] = tex2D<T>(cur, xo, yo);
 116       return;
 117     }
 118
 119     // Calculate spatial prediction
 120     T a = tex2D<T>(cur, xo - 3, yo - 1);
 121     T b = tex2D<T>(cur, xo - 2, yo - 1);
 122     T c = tex2D<T>(cur, xo - 1, yo - 1);
 123     T d = tex2D<T>(cur, xo - 0, yo - 1);
 124     T e = tex2D<T>(cur, xo + 1, yo - 1);
 125     T f = tex2D<T>(cur, xo + 2, yo - 1);
 126     T g = tex2D<T>(cur, xo + 3, yo - 1);
 127
 128     T h = tex2D<T>(cur, xo - 3, yo + 1);
 129     T i = tex2D<T>(cur, xo - 2, yo + 1);
 130     T j = tex2D<T>(cur, xo - 1, yo + 1);
 131     T k = tex2D<T>(cur, xo - 0, yo + 1);
 132     T l = tex2D<T>(cur, xo + 1, yo + 1);
 133     T m = tex2D<T>(cur, xo + 2, yo + 1);
 134     T n = tex2D<T>(cur, xo + 3, yo + 1);
 135
 136     T spatial_pred =
 137         spatial_predictor(a, b, c, d, e, f, g, h, i, j, k, l, m, n);
 138
 139     // Calculate temporal prediction
 140     int is_second_field = !(parity ^ tff);
 141
 142     cudaTextureObject_t prev2 = prev;
 143     cudaTextureObject_t prev1 = is_second_field ? cur : prev;
 144     cudaTextureObject_t next1 = is_second_field ? next : cur;
 145     cudaTextureObject_t next2 = next;
 146
 147     T A = tex2D<T>(prev2, xo,  yo - 1);
 148     T B = tex2D<T>(prev2, xo,  yo + 1);
 149     T C = tex2D<T>(prev1, xo,  yo - 2);
 150     T D = tex2D<T>(prev1, xo,  yo + 0);
 151     T E = tex2D<T>(prev1, xo,  yo + 2);
 152     T F = tex2D<T>(cur,   xo,  yo - 1);
 153     T G = tex2D<T>(cur,   xo,  yo + 1);
 154     T H = tex2D<T>(next1, xo,  yo - 2);
 155     T I = tex2D<T>(next1, xo,  yo + 0);
 156     T J = tex2D<T>(next1, xo,  yo + 2);
 157     T K = tex2D<T>(next2, xo,  yo - 1);
 158     T L = tex2D<T>(next2, xo,  yo + 1);
 159
 160     spatial_pred = temporal_predictor(A, B, C, D, E, F, G, H, I, J, K, L,
 161                                       spatial_pred, skip_spatial_check);
 162
 163     dst[yo*dst_pitch+xo] = spatial_pred;
 164 }
 165
 166 template <typename T>
 167 __inline__ __device__ void yadif_double(T *dst,
 168                                         cudaTextureObject_t prev,
 169                                         cudaTextureObject_t cur,
 170                                         cudaTextureObject_t next,
 171                                         int dst_width, int dst_height, int dst_pitch,
 172                                         int src_width, int src_height,
 173                                         int parity, int tff, bool skip_spatial_check)
 174 {
 175     int xo = blockIdx.x * blockDim.x + threadIdx.x;
 176     int yo = blockIdx.y * blockDim.y + threadIdx.y;
 177
 178     if (xo >= dst_width || yo >= dst_height) {
 179         return;
 180     }
 181
 182     if (yo % 2 == parity) {
 183       // Don't modify the primary field
 184       dst[yo*dst_pitch+xo] = tex2D<T>(cur, xo, yo);
 185       return;
 186     }
 187
 188     T a = tex2D<T>(cur, xo - 3, yo - 1);
 189     T b = tex2D<T>(cur, xo - 2, yo - 1);
 190     T c = tex2D<T>(cur, xo - 1, yo - 1);
 191     T d = tex2D<T>(cur, xo - 0, yo - 1);
 192     T e = tex2D<T>(cur, xo + 1, yo - 1);
 193     T f = tex2D<T>(cur, xo + 2, yo - 1);
 194     T g = tex2D<T>(cur, xo + 3, yo - 1);
 195
 196     T h = tex2D<T>(cur, xo - 3, yo + 1);
 197     T i = tex2D<T>(cur, xo - 2, yo + 1);
 198     T j = tex2D<T>(cur, xo - 1, yo + 1);
 199     T k = tex2D<T>(cur, xo - 0, yo + 1);
 200     T l = tex2D<T>(cur, xo + 1, yo + 1);
 201     T m = tex2D<T>(cur, xo + 2, yo + 1);
 202     T n = tex2D<T>(cur, xo + 3, yo + 1);
 203
 204     T spatial_pred;
 205     spatial_pred.x =
 206         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);
 207     spatial_pred.y =
 208         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);
 209
 210     // Calculate temporal prediction
 211     int is_second_field = !(parity ^ tff);
 212
 213     cudaTextureObject_t prev2 = prev;
 214     cudaTextureObject_t prev1 = is_second_field ? cur : prev;
 215     cudaTextureObject_t next1 = is_second_field ? next : cur;
 216     cudaTextureObject_t next2 = next;
 217
 218     T A = tex2D<T>(prev2, xo,  yo - 1);
 219     T B = tex2D<T>(prev2, xo,  yo + 1);
 220     T C = tex2D<T>(prev1, xo,  yo - 2);
 221     T D = tex2D<T>(prev1, xo,  yo + 0);
 222     T E = tex2D<T>(prev1, xo,  yo + 2);
 223     T F = tex2D<T>(cur,   xo,  yo - 1);
 224     T G = tex2D<T>(cur,   xo,  yo + 1);
 225     T H = tex2D<T>(next1, xo,  yo - 2);
 226     T I = tex2D<T>(next1, xo,  yo + 0);
 227     T J = tex2D<T>(next1, xo,  yo + 2);
 228     T K = tex2D<T>(next2, xo,  yo - 1);
 229     T L = tex2D<T>(next2, xo,  yo + 1);
 230
 231     spatial_pred.x =
 232         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,
 233                            spatial_pred.x, skip_spatial_check);
 234     spatial_pred.y =
 235         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,
 236                            spatial_pred.y, skip_spatial_check);
 237
 238     dst[yo*dst_pitch+xo] = spatial_pred;
 239 }
 240
 241 extern "C" {
 242
 243 __global__ void yadif_uchar(unsigned char *dst,
 244                             cudaTextureObject_t prev,
 245                             cudaTextureObject_t cur,
 246                             cudaTextureObject_t next,
 247                             int dst_width, int dst_height, int dst_pitch,
 248                             int src_width, int src_height,
 249                             int parity, int tff, bool skip_spatial_check)
 250 {
 251     yadif_single(dst, prev, cur, next,
 252                  dst_width, dst_height, dst_pitch,
 253                  src_width, src_height,
 254                  parity, tff, skip_spatial_check);
 255 }
 256
 257 __global__ void yadif_ushort(unsigned short *dst,
 258                             cudaTextureObject_t prev,
 259                             cudaTextureObject_t cur,
 260                             cudaTextureObject_t next,
 261                             int dst_width, int dst_height, int dst_pitch,
 262                             int src_width, int src_height,
 263                             int parity, int tff, bool skip_spatial_check)
 264 {
 265     yadif_single(dst, prev, cur, next,
 266                  dst_width, dst_height, dst_pitch,
 267                  src_width, src_height,
 268                  parity, tff, skip_spatial_check);
 269 }
 270
 271 __global__ void yadif_uchar2(uchar2 *dst,
 272                             cudaTextureObject_t prev,
 273                             cudaTextureObject_t cur,
 274                             cudaTextureObject_t next,
 275                             int dst_width, int dst_height, int dst_pitch,
 276                             int src_width, int src_height,
 277                             int parity, int tff, bool skip_spatial_check)
 278 {
 279     yadif_double(dst, prev, cur, next,
 280                  dst_width, dst_height, dst_pitch,
 281                  src_width, src_height,
 282                  parity, tff, skip_spatial_check);
 283 }
 284
 285 __global__ void yadif_ushort2(ushort2 *dst,
 286                             cudaTextureObject_t prev,
 287                             cudaTextureObject_t cur,
 288                             cudaTextureObject_t next,
 289                             int dst_width, int dst_height, int dst_pitch,
 290                             int src_width, int src_height,
 291                             int parity, int tff, bool skip_spatial_check)
 292 {
 293     yadif_double(dst, prev, cur, next,
 294                  dst_width, dst_height, dst_pitch,
 295                  src_width, src_height,
 296                  parity, tff, skip_spatial_check);
 297 }
 298
 299 } /* extern "C" */