2 * Copyright (C) 2018 Philip Langdale <philipl@overt.org>
4 * This file is part of FFmpeg.
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.
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.
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
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)
25 int spatial_pred = (d + k)/2;
26 int spatial_score = abs(c - j) + abs(d - k) + abs(e - l);
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;
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;
51 __inline__ __device__ int max3(int a, int b, int c)
57 __inline__ __device__ int min3(int a, int b, int c)
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)
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;
78 int diff = max3(tdiff0, tdiff1, tdiff2);
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);
86 if (spatial_pred > p2 + diff) {
87 spatial_pred = p2 + diff;
89 if (spatial_pred < p2 - diff) {
90 spatial_pred = p2 - diff;
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)
106 int xo = blockIdx.x * blockDim.x + threadIdx.x;
107 int yo = blockIdx.y * blockDim.y + threadIdx.y;
109 if (xo >= dst_width || yo >= dst_height) {
113 // Don't modify the primary field
114 if (yo % 2 == parity) {
115 dst[yo*dst_pitch+xo] = tex2D<T>(cur, xo, yo);
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);
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);
137 spatial_predictor(a, b, c, d, e, f, g, h, i, j, k, l, m, n);
139 // Calculate temporal prediction
140 int is_second_field = !(parity ^ tff);
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;
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);
160 spatial_pred = temporal_predictor(A, B, C, D, E, F, G, H, I, J, K, L,
161 spatial_pred, skip_spatial_check);
163 dst[yo*dst_pitch+xo] = spatial_pred;
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)
175 int xo = blockIdx.x * blockDim.x + threadIdx.x;
176 int yo = blockIdx.y * blockDim.y + threadIdx.y;
178 if (xo >= dst_width || yo >= dst_height) {
182 if (yo % 2 == parity) {
183 // Don't modify the primary field
184 dst[yo*dst_pitch+xo] = tex2D<T>(cur, xo, yo);
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);
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);
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);
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);
210 // Calculate temporal prediction
211 int is_second_field = !(parity ^ tff);
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;
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);
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);
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);
238 dst[yo*dst_pitch+xo] = spatial_pred;
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)
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);
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)
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);
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)
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);
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)
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);