65a902c66b23406b662707aa76fef10fc24d515c
[ffmpeg.git] / 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_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),
206         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) };
207
208     // Calculate temporal prediction
209     int is_second_field = !(parity ^ tff);
210
211     cudaTextureObject_t prev2 = prev;
212     cudaTextureObject_t prev1 = is_second_field ? cur : prev;
213     cudaTextureObject_t next1 = is_second_field ? next : cur;
214     cudaTextureObject_t next2 = next;
215
216     T A = tex2D<T>(prev2, xo,  yo - 1);
217     T B = tex2D<T>(prev2, xo,  yo + 1);
218     T C = tex2D<T>(prev1, xo,  yo - 2);
219     T D = tex2D<T>(prev1, xo,  yo + 0);
220     T E = tex2D<T>(prev1, xo,  yo + 2);
221     T F = tex2D<T>(cur,   xo,  yo - 1);
222     T G = tex2D<T>(cur,   xo,  yo + 1);
223     T H = tex2D<T>(next1, xo,  yo - 2);
224     T I = tex2D<T>(next1, xo,  yo + 0);
225     T J = tex2D<T>(next1, xo,  yo + 2);
226     T K = tex2D<T>(next2, xo,  yo - 1);
227     T L = tex2D<T>(next2, xo,  yo + 1);
228
229     spatial_pred = {
230         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,
231                            spatial_pred.x, skip_spatial_check),
232         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,
233                            spatial_pred.y, skip_spatial_check) };
234
235     dst[yo*dst_pitch+xo] = spatial_pred;
236 }
237
238 extern "C" {
239
240 __global__ void yadif_uchar(unsigned char *dst,
241                             cudaTextureObject_t prev,
242                             cudaTextureObject_t cur,
243                             cudaTextureObject_t next,
244                             int dst_width, int dst_height, int dst_pitch,
245                             int src_width, int src_height,
246                             int parity, int tff, bool skip_spatial_check)
247 {
248     yadif_single(dst, prev, cur, next,
249                  dst_width, dst_height, dst_pitch,
250                  src_width, src_height,
251                  parity, tff, skip_spatial_check);
252 }
253
254 __global__ void yadif_ushort(unsigned short *dst,
255                             cudaTextureObject_t prev,
256                             cudaTextureObject_t cur,
257                             cudaTextureObject_t next,
258                             int dst_width, int dst_height, int dst_pitch,
259                             int src_width, int src_height,
260                             int parity, int tff, bool skip_spatial_check)
261 {
262     yadif_single(dst, prev, cur, next,
263                  dst_width, dst_height, dst_pitch,
264                  src_width, src_height,
265                  parity, tff, skip_spatial_check);
266 }
267
268 __global__ void yadif_uchar2(uchar2 *dst,
269                             cudaTextureObject_t prev,
270                             cudaTextureObject_t cur,
271                             cudaTextureObject_t next,
272                             int dst_width, int dst_height, int dst_pitch,
273                             int src_width, int src_height,
274                             int parity, int tff, bool skip_spatial_check)
275 {
276     yadif_double(dst, prev, cur, next,
277                  dst_width, dst_height, dst_pitch,
278                  src_width, src_height,
279                  parity, tff, skip_spatial_check);
280 }
281
282 __global__ void yadif_ushort2(ushort2 *dst,
283                             cudaTextureObject_t prev,
284                             cudaTextureObject_t cur,
285                             cudaTextureObject_t next,
286                             int dst_width, int dst_height, int dst_pitch,
287                             int src_width, int src_height,
288                             int parity, int tff, bool skip_spatial_check)
289 {
290     yadif_double(dst, prev, cur, next,
291                  dst_width, dst_height, dst_pitch,
292                  src_width, src_height,
293                  parity, tff, skip_spatial_check);
294 }
295
296 } /* extern "C" */