85e1aac5ebdebf3218510fc0e60a814cfeec031f
[ffmpeg.git] / libavfilter / vf_yadif_cuda.c
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 #include <cuda.h>
22 #include "libavutil/avassert.h"
23 #include "libavutil/hwcontext_cuda.h"
24 #include "libavutil/cuda_check.h"
25 #include "internal.h"
26 #include "yadif.h"
27
28 extern char vf_yadif_cuda_ptx[];
29
30 typedef struct DeintCUDAContext {
31     YADIFContext yadif;
32
33     AVCUDADeviceContext *hwctx;
34     AVBufferRef         *device_ref;
35     AVBufferRef         *input_frames_ref;
36     AVHWFramesContext   *input_frames;
37
38     CUcontext   cu_ctx;
39     CUstream    stream;
40     CUmodule    cu_module;
41     CUfunction  cu_func_uchar;
42     CUfunction  cu_func_uchar2;
43     CUfunction  cu_func_ushort;
44     CUfunction  cu_func_ushort2;
45 } DeintCUDAContext;
46
47 #define DIV_UP(a, b) ( ((a) + (b) - 1) / (b) )
48 #define ALIGN_UP(a, b) (((a) + (b) - 1) & ~((b) - 1))
49 #define BLOCKX 32
50 #define BLOCKY 16
51
52 #define CHECK_CU(x) FF_CUDA_CHECK(ctx, x)
53
54 static CUresult call_kernel(AVFilterContext *ctx, CUfunction func,
55                             CUdeviceptr prev, CUdeviceptr cur, CUdeviceptr next,
56                             CUarray_format format, int channels,
57                             int src_width,  // Width is pixels per channel
58                             int src_height, // Height is pixels per channel
59                             int src_pitch,  // Pitch is bytes
60                             CUdeviceptr dst,
61                             int dst_width,  // Width is pixels per channel
62                             int dst_height, // Height is pixels per channel
63                             int dst_pitch,  // Pitch is pixels per channel
64                             int parity, int tff)
65 {
66     DeintCUDAContext *s = ctx->priv;
67     CUtexObject tex_prev = 0, tex_cur = 0, tex_next = 0;
68     int ret;
69     int skip_spatial_check = s->yadif.mode&2;
70
71     void *args[] = { &dst, &tex_prev, &tex_cur, &tex_next,
72                      &dst_width, &dst_height, &dst_pitch,
73                      &src_width, &src_height, &parity, &tff,
74                      &skip_spatial_check };
75
76     CUDA_TEXTURE_DESC tex_desc = {
77         .filterMode = CU_TR_FILTER_MODE_POINT,
78         .flags = CU_TRSF_READ_AS_INTEGER,
79     };
80
81     CUDA_RESOURCE_DESC res_desc = {
82         .resType = CU_RESOURCE_TYPE_PITCH2D,
83         .res.pitch2D.format = format,
84         .res.pitch2D.numChannels = channels,
85         .res.pitch2D.width = src_width,
86         .res.pitch2D.height = src_height,
87         .res.pitch2D.pitchInBytes = src_pitch,
88     };
89
90     res_desc.res.pitch2D.devPtr = (CUdeviceptr)prev;
91     ret = CHECK_CU(cuTexObjectCreate(&tex_prev, &res_desc, &tex_desc, NULL));
92     if (ret < 0)
93         goto exit;
94
95     res_desc.res.pitch2D.devPtr = (CUdeviceptr)cur;
96     ret = CHECK_CU(cuTexObjectCreate(&tex_cur, &res_desc, &tex_desc, NULL));
97     if (ret < 0)
98         goto exit;
99
100     res_desc.res.pitch2D.devPtr = (CUdeviceptr)next;
101     ret = CHECK_CU(cuTexObjectCreate(&tex_next, &res_desc, &tex_desc, NULL));
102     if (ret < 0)
103         goto exit;
104
105     ret = CHECK_CU(cuLaunchKernel(func,
106                                   DIV_UP(dst_width, BLOCKX), DIV_UP(dst_height, BLOCKY), 1,
107                                   BLOCKX, BLOCKY, 1,
108                                   0, s->stream, args, NULL));
109
110 exit:
111     if (tex_prev)
112         CHECK_CU(cuTexObjectDestroy(tex_prev));
113     if (tex_cur)
114         CHECK_CU(cuTexObjectDestroy(tex_cur));
115     if (tex_next)
116         CHECK_CU(cuTexObjectDestroy(tex_next));
117
118     return ret;
119 }
120
121 static void filter(AVFilterContext *ctx, AVFrame *dst,
122                    int parity, int tff)
123 {
124     DeintCUDAContext *s = ctx->priv;
125     YADIFContext *y = &s->yadif;
126     CUcontext dummy;
127     int i, ret;
128
129     ret = CHECK_CU(cuCtxPushCurrent(s->cu_ctx));
130     if (ret < 0)
131         return;
132
133     for (i = 0; i < y->csp->nb_components; i++) {
134         CUfunction func;
135         CUarray_format format;
136         int pixel_size, channels;
137         const AVComponentDescriptor *comp = &y->csp->comp[i];
138
139         if (comp->plane < i) {
140             // We process planes as a whole, so don't reprocess
141             // them for additional components
142             continue;
143         }
144
145         pixel_size = (comp->depth + comp->shift) / 8;
146         channels = comp->step / pixel_size;
147         if (pixel_size > 2 || channels > 2) {
148             av_log(ctx, AV_LOG_ERROR, "Unsupported pixel format: %s\n", y->csp->name);
149             goto exit;
150         }
151         switch (pixel_size) {
152         case 1:
153             func = channels == 1 ? s->cu_func_uchar : s->cu_func_uchar2;
154             format = CU_AD_FORMAT_UNSIGNED_INT8;
155             break;
156         case 2:
157             func = channels == 1 ? s->cu_func_ushort : s->cu_func_ushort2;
158             format = CU_AD_FORMAT_UNSIGNED_INT16;
159             break;
160         default:
161             av_log(ctx, AV_LOG_ERROR, "Unsupported pixel format: %s\n", y->csp->name);
162             goto exit;
163         }
164         av_log(ctx, AV_LOG_TRACE,
165                "Deinterlacing plane %d: pixel_size: %d channels: %d\n",
166                comp->plane, pixel_size, channels);
167         call_kernel(ctx, func,
168                     (CUdeviceptr)y->prev->data[i],
169                     (CUdeviceptr)y->cur->data[i],
170                     (CUdeviceptr)y->next->data[i],
171                     format, channels,
172                     AV_CEIL_RSHIFT(y->cur->width, i ? y->csp->log2_chroma_w : 0),
173                     AV_CEIL_RSHIFT(y->cur->height, i ? y->csp->log2_chroma_h : 0),
174                     y->cur->linesize[i],
175                     (CUdeviceptr)dst->data[i],
176                     AV_CEIL_RSHIFT(dst->width, i ? y->csp->log2_chroma_w : 0),
177                     AV_CEIL_RSHIFT(dst->height, i ? y->csp->log2_chroma_h : 0),
178                     dst->linesize[i] / comp->step,
179                     parity, tff);
180     }
181
182     CHECK_CU(cuStreamSynchronize(s->stream));
183
184 exit:
185     CHECK_CU(cuCtxPopCurrent(&dummy));
186     return;
187 }
188
189 static av_cold void deint_cuda_uninit(AVFilterContext *ctx)
190 {
191     CUcontext dummy;
192     DeintCUDAContext *s = ctx->priv;
193     YADIFContext *y = &s->yadif;
194
195     if (s->cu_module) {
196         CHECK_CU(cuCtxPushCurrent(s->cu_ctx));
197         CHECK_CU(cuModuleUnload(s->cu_module));
198         CHECK_CU(cuCtxPopCurrent(&dummy));
199     }
200
201     av_frame_free(&y->prev);
202     av_frame_free(&y->cur);
203     av_frame_free(&y->next);
204
205     av_buffer_unref(&s->device_ref);
206     s->hwctx = NULL;
207     av_buffer_unref(&s->input_frames_ref);
208     s->input_frames = NULL;
209 }
210
211 static int deint_cuda_query_formats(AVFilterContext *ctx)
212 {
213     enum AVPixelFormat pix_fmts[] = {
214         AV_PIX_FMT_CUDA, AV_PIX_FMT_NONE,
215     };
216     int ret;
217
218     if ((ret = ff_formats_ref(ff_make_format_list(pix_fmts),
219                               &ctx->inputs[0]->out_formats)) < 0)
220         return ret;
221     if ((ret = ff_formats_ref(ff_make_format_list(pix_fmts),
222                               &ctx->outputs[0]->in_formats)) < 0)
223         return ret;
224
225     return 0;
226 }
227
228 static int config_input(AVFilterLink *inlink)
229 {
230     AVFilterContext *ctx = inlink->dst;
231     DeintCUDAContext *s  = ctx->priv;
232
233     if (!inlink->hw_frames_ctx) {
234         av_log(ctx, AV_LOG_ERROR, "A hardware frames reference is "
235                "required to associate the processing device.\n");
236         return AVERROR(EINVAL);
237     }
238
239     s->input_frames_ref = av_buffer_ref(inlink->hw_frames_ctx);
240     if (!s->input_frames_ref) {
241         av_log(ctx, AV_LOG_ERROR, "A input frames reference create "
242                "failed.\n");
243         return AVERROR(ENOMEM);
244     }
245     s->input_frames = (AVHWFramesContext*)s->input_frames_ref->data;
246
247     return 0;
248 }
249
250 static int config_output(AVFilterLink *link)
251 {
252     AVHWFramesContext *output_frames;
253     AVFilterContext *ctx = link->src;
254     DeintCUDAContext *s = ctx->priv;
255     YADIFContext *y = &s->yadif;
256     int ret = 0;
257     CUcontext dummy;
258
259     av_assert0(s->input_frames);
260     s->device_ref = av_buffer_ref(s->input_frames->device_ref);
261     if (!s->device_ref) {
262         av_log(ctx, AV_LOG_ERROR, "A device reference create "
263                "failed.\n");
264         return AVERROR(ENOMEM);
265     }
266     s->hwctx = ((AVHWDeviceContext*)s->device_ref->data)->hwctx;
267     s->cu_ctx = s->hwctx->cuda_ctx;
268     s->stream = s->hwctx->stream;
269
270     link->hw_frames_ctx = av_hwframe_ctx_alloc(s->device_ref);
271     if (!link->hw_frames_ctx) {
272         av_log(ctx, AV_LOG_ERROR, "Failed to create HW frame context "
273                "for output.\n");
274         ret = AVERROR(ENOMEM);
275         goto exit;
276     }
277
278     output_frames = (AVHWFramesContext*)link->hw_frames_ctx->data;
279
280     output_frames->format    = AV_PIX_FMT_CUDA;
281     output_frames->sw_format = s->input_frames->sw_format;
282     output_frames->width     = ctx->inputs[0]->w;
283     output_frames->height    = ctx->inputs[0]->h;
284
285     output_frames->initial_pool_size = 4;
286
287     ret = ff_filter_init_hw_frames(ctx, link, 10);
288     if (ret < 0)
289         goto exit;
290
291     ret = av_hwframe_ctx_init(link->hw_frames_ctx);
292     if (ret < 0) {
293         av_log(ctx, AV_LOG_ERROR, "Failed to initialise CUDA frame "
294                "context for output: %d\n", ret);
295         goto exit;
296     }
297
298     link->time_base.num = ctx->inputs[0]->time_base.num;
299     link->time_base.den = ctx->inputs[0]->time_base.den * 2;
300     link->w             = ctx->inputs[0]->w;
301     link->h             = ctx->inputs[0]->h;
302
303     if(y->mode & 1)
304         link->frame_rate = av_mul_q(ctx->inputs[0]->frame_rate,
305                                     (AVRational){2, 1});
306
307     if (link->w < 3 || link->h < 3) {
308         av_log(ctx, AV_LOG_ERROR, "Video of less than 3 columns or lines is not supported\n");
309         ret = AVERROR(EINVAL);
310         goto exit;
311     }
312
313     y->csp = av_pix_fmt_desc_get(output_frames->sw_format);
314     y->filter = filter;
315
316     ret = CHECK_CU(cuCtxPushCurrent(s->cu_ctx));
317     if (ret < 0)
318         goto exit;
319
320     ret = CHECK_CU(cuModuleLoadData(&s->cu_module, vf_yadif_cuda_ptx));
321     if (ret < 0)
322         goto exit;
323
324     ret = CHECK_CU(cuModuleGetFunction(&s->cu_func_uchar, s->cu_module, "yadif_uchar"));
325     if (ret < 0)
326         goto exit;
327
328     ret = CHECK_CU(cuModuleGetFunction(&s->cu_func_uchar2, s->cu_module, "yadif_uchar2"));
329     if (ret < 0)
330         goto exit;
331
332     ret = CHECK_CU(cuModuleGetFunction(&s->cu_func_ushort, s->cu_module, "yadif_ushort"));
333     if (ret < 0)
334         goto exit;
335
336     ret = CHECK_CU(cuModuleGetFunction(&s->cu_func_ushort2, s->cu_module, "yadif_ushort2"));
337     if (ret < 0)
338         goto exit;
339
340 exit:
341     CHECK_CU(cuCtxPopCurrent(&dummy));
342
343     return ret;
344 }
345
346 static const AVClass yadif_cuda_class = {
347     .class_name = "yadif_cuda",
348     .item_name  = av_default_item_name,
349     .option     = ff_yadif_options,
350     .version    = LIBAVUTIL_VERSION_INT,
351     .category   = AV_CLASS_CATEGORY_FILTER,
352 };
353
354 static const AVFilterPad deint_cuda_inputs[] = {
355     {
356         .name          = "default",
357         .type          = AVMEDIA_TYPE_VIDEO,
358         .filter_frame  = ff_yadif_filter_frame,
359         .config_props  = config_input,
360     },
361     { NULL }
362 };
363
364 static const AVFilterPad deint_cuda_outputs[] = {
365     {
366         .name          = "default",
367         .type          = AVMEDIA_TYPE_VIDEO,
368         .request_frame = ff_yadif_request_frame,
369         .config_props  = config_output,
370     },
371     { NULL }
372 };
373
374 AVFilter ff_vf_yadif_cuda = {
375     .name           = "yadif_cuda",
376     .description    = NULL_IF_CONFIG_SMALL("Deinterlace CUDA frames"),
377     .priv_size      = sizeof(DeintCUDAContext),
378     .priv_class     = &yadif_cuda_class,
379     .uninit         = deint_cuda_uninit,
380     .query_formats  = deint_cuda_query_formats,
381     .inputs         = deint_cuda_inputs,
382     .outputs        = deint_cuda_outputs,
383     .flags          = AVFILTER_FLAG_SUPPORT_TIMELINE_INTERNAL,
384     .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
385 };