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
21 #include "libavutil/avassert.h"
22 #include "libavutil/hwcontext_cuda_internal.h"
23 #include "libavutil/cuda_check.h"
27 extern char vf_yadif_cuda_ptx[];
29 typedef struct DeintCUDAContext {
32 AVCUDADeviceContext *hwctx;
33 AVBufferRef *device_ref;
34 AVBufferRef *input_frames_ref;
35 AVHWFramesContext *input_frames;
40 CUfunction cu_func_uchar;
41 CUfunction cu_func_uchar2;
42 CUfunction cu_func_ushort;
43 CUfunction cu_func_ushort2;
46 #define DIV_UP(a, b) ( ((a) + (b) - 1) / (b) )
47 #define ALIGN_UP(a, b) (((a) + (b) - 1) & ~((b) - 1))
51 #define CHECK_CU(x) FF_CUDA_CHECK_DL(ctx, s->hwctx->internal->cuda_dl, x)
53 static CUresult call_kernel(AVFilterContext *ctx, CUfunction func,
54 CUdeviceptr prev, CUdeviceptr cur, CUdeviceptr next,
55 CUarray_format format, int channels,
56 int src_width, // Width is pixels per channel
57 int src_height, // Height is pixels per channel
58 int src_pitch, // Pitch is bytes
60 int dst_width, // Width is pixels per channel
61 int dst_height, // Height is pixels per channel
62 int dst_pitch, // Pitch is pixels per channel
65 DeintCUDAContext *s = ctx->priv;
66 CudaFunctions *cu = s->hwctx->internal->cuda_dl;
67 CUtexObject tex_prev = 0, tex_cur = 0, tex_next = 0;
69 int skip_spatial_check = s->yadif.mode&2;
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 };
76 CUDA_TEXTURE_DESC tex_desc = {
77 .filterMode = CU_TR_FILTER_MODE_POINT,
78 .flags = CU_TRSF_READ_AS_INTEGER,
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,
90 res_desc.res.pitch2D.devPtr = (CUdeviceptr)prev;
91 ret = CHECK_CU(cu->cuTexObjectCreate(&tex_prev, &res_desc, &tex_desc, NULL));
95 res_desc.res.pitch2D.devPtr = (CUdeviceptr)cur;
96 ret = CHECK_CU(cu->cuTexObjectCreate(&tex_cur, &res_desc, &tex_desc, NULL));
100 res_desc.res.pitch2D.devPtr = (CUdeviceptr)next;
101 ret = CHECK_CU(cu->cuTexObjectCreate(&tex_next, &res_desc, &tex_desc, NULL));
105 ret = CHECK_CU(cu->cuLaunchKernel(func,
106 DIV_UP(dst_width, BLOCKX), DIV_UP(dst_height, BLOCKY), 1,
108 0, s->stream, args, NULL));
112 CHECK_CU(cu->cuTexObjectDestroy(tex_prev));
114 CHECK_CU(cu->cuTexObjectDestroy(tex_cur));
116 CHECK_CU(cu->cuTexObjectDestroy(tex_next));
121 static void filter(AVFilterContext *ctx, AVFrame *dst,
124 DeintCUDAContext *s = ctx->priv;
125 YADIFContext *y = &s->yadif;
126 CudaFunctions *cu = s->hwctx->internal->cuda_dl;
130 ret = CHECK_CU(cu->cuCtxPushCurrent(s->cu_ctx));
134 for (i = 0; i < y->csp->nb_components; i++) {
136 CUarray_format format;
137 int pixel_size, channels;
138 const AVComponentDescriptor *comp = &y->csp->comp[i];
140 if (comp->plane < i) {
141 // We process planes as a whole, so don't reprocess
142 // them for additional components
146 pixel_size = (comp->depth + comp->shift) / 8;
147 channels = comp->step / pixel_size;
148 if (pixel_size > 2 || channels > 2) {
149 av_log(ctx, AV_LOG_ERROR, "Unsupported pixel format: %s\n", y->csp->name);
152 switch (pixel_size) {
154 func = channels == 1 ? s->cu_func_uchar : s->cu_func_uchar2;
155 format = CU_AD_FORMAT_UNSIGNED_INT8;
158 func = channels == 1 ? s->cu_func_ushort : s->cu_func_ushort2;
159 format = CU_AD_FORMAT_UNSIGNED_INT16;
162 av_log(ctx, AV_LOG_ERROR, "Unsupported pixel format: %s\n", y->csp->name);
165 av_log(ctx, AV_LOG_TRACE,
166 "Deinterlacing plane %d: pixel_size: %d channels: %d\n",
167 comp->plane, pixel_size, channels);
168 call_kernel(ctx, func,
169 (CUdeviceptr)y->prev->data[i],
170 (CUdeviceptr)y->cur->data[i],
171 (CUdeviceptr)y->next->data[i],
173 AV_CEIL_RSHIFT(y->cur->width, i ? y->csp->log2_chroma_w : 0),
174 AV_CEIL_RSHIFT(y->cur->height, i ? y->csp->log2_chroma_h : 0),
176 (CUdeviceptr)dst->data[i],
177 AV_CEIL_RSHIFT(dst->width, i ? y->csp->log2_chroma_w : 0),
178 AV_CEIL_RSHIFT(dst->height, i ? y->csp->log2_chroma_h : 0),
179 dst->linesize[i] / comp->step,
184 CHECK_CU(cu->cuCtxPopCurrent(&dummy));
188 static av_cold void deint_cuda_uninit(AVFilterContext *ctx)
191 DeintCUDAContext *s = ctx->priv;
192 YADIFContext *y = &s->yadif;
194 if (s->hwctx && s->cu_module) {
195 CudaFunctions *cu = s->hwctx->internal->cuda_dl;
196 CHECK_CU(cu->cuCtxPushCurrent(s->cu_ctx));
197 CHECK_CU(cu->cuModuleUnload(s->cu_module));
198 CHECK_CU(cu->cuCtxPopCurrent(&dummy));
201 av_frame_free(&y->prev);
202 av_frame_free(&y->cur);
203 av_frame_free(&y->next);
205 av_buffer_unref(&s->device_ref);
207 av_buffer_unref(&s->input_frames_ref);
208 s->input_frames = NULL;
211 static int deint_cuda_query_formats(AVFilterContext *ctx)
213 enum AVPixelFormat pix_fmts[] = {
214 AV_PIX_FMT_CUDA, AV_PIX_FMT_NONE,
218 if ((ret = ff_formats_ref(ff_make_format_list(pix_fmts),
219 &ctx->inputs[0]->outcfg.formats)) < 0)
221 if ((ret = ff_formats_ref(ff_make_format_list(pix_fmts),
222 &ctx->outputs[0]->incfg.formats)) < 0)
228 static int config_input(AVFilterLink *inlink)
230 AVFilterContext *ctx = inlink->dst;
231 DeintCUDAContext *s = ctx->priv;
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);
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 "
243 return AVERROR(ENOMEM);
245 s->input_frames = (AVHWFramesContext*)s->input_frames_ref->data;
250 static int config_output(AVFilterLink *link)
252 AVHWFramesContext *output_frames;
253 AVFilterContext *ctx = link->src;
254 DeintCUDAContext *s = ctx->priv;
255 YADIFContext *y = &s->yadif;
260 av_assert0(s->input_frames);
261 s->device_ref = av_buffer_ref(s->input_frames->device_ref);
262 if (!s->device_ref) {
263 av_log(ctx, AV_LOG_ERROR, "A device reference create "
265 return AVERROR(ENOMEM);
267 s->hwctx = ((AVHWDeviceContext*)s->device_ref->data)->hwctx;
268 s->cu_ctx = s->hwctx->cuda_ctx;
269 s->stream = s->hwctx->stream;
270 cu = s->hwctx->internal->cuda_dl;
272 link->hw_frames_ctx = av_hwframe_ctx_alloc(s->device_ref);
273 if (!link->hw_frames_ctx) {
274 av_log(ctx, AV_LOG_ERROR, "Failed to create HW frame context "
276 ret = AVERROR(ENOMEM);
280 output_frames = (AVHWFramesContext*)link->hw_frames_ctx->data;
282 output_frames->format = AV_PIX_FMT_CUDA;
283 output_frames->sw_format = s->input_frames->sw_format;
284 output_frames->width = ctx->inputs[0]->w;
285 output_frames->height = ctx->inputs[0]->h;
287 output_frames->initial_pool_size = 4;
289 ret = ff_filter_init_hw_frames(ctx, link, 10);
293 ret = av_hwframe_ctx_init(link->hw_frames_ctx);
295 av_log(ctx, AV_LOG_ERROR, "Failed to initialise CUDA frame "
296 "context for output: %d\n", ret);
300 link->time_base.num = ctx->inputs[0]->time_base.num;
301 link->time_base.den = ctx->inputs[0]->time_base.den * 2;
302 link->w = ctx->inputs[0]->w;
303 link->h = ctx->inputs[0]->h;
306 link->frame_rate = av_mul_q(ctx->inputs[0]->frame_rate,
309 if (link->w < 3 || link->h < 3) {
310 av_log(ctx, AV_LOG_ERROR, "Video of less than 3 columns or lines is not supported\n");
311 ret = AVERROR(EINVAL);
315 y->csp = av_pix_fmt_desc_get(output_frames->sw_format);
318 ret = CHECK_CU(cu->cuCtxPushCurrent(s->cu_ctx));
322 ret = CHECK_CU(cu->cuModuleLoadData(&s->cu_module, vf_yadif_cuda_ptx));
326 ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar, s->cu_module, "yadif_uchar"));
330 ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar2, s->cu_module, "yadif_uchar2"));
334 ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_ushort, s->cu_module, "yadif_ushort"));
338 ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_ushort2, s->cu_module, "yadif_ushort2"));
343 CHECK_CU(cu->cuCtxPopCurrent(&dummy));
348 static const AVClass yadif_cuda_class = {
349 .class_name = "yadif_cuda",
350 .item_name = av_default_item_name,
351 .option = ff_yadif_options,
352 .version = LIBAVUTIL_VERSION_INT,
353 .category = AV_CLASS_CATEGORY_FILTER,
356 static const AVFilterPad deint_cuda_inputs[] = {
359 .type = AVMEDIA_TYPE_VIDEO,
360 .filter_frame = ff_yadif_filter_frame,
361 .config_props = config_input,
366 static const AVFilterPad deint_cuda_outputs[] = {
369 .type = AVMEDIA_TYPE_VIDEO,
370 .request_frame = ff_yadif_request_frame,
371 .config_props = config_output,
376 AVFilter ff_vf_yadif_cuda = {
377 .name = "yadif_cuda",
378 .description = NULL_IF_CONFIG_SMALL("Deinterlace CUDA frames"),
379 .priv_size = sizeof(DeintCUDAContext),
380 .priv_class = &yadif_cuda_class,
381 .uninit = deint_cuda_uninit,
382 .query_formats = deint_cuda_query_formats,
383 .inputs = deint_cuda_inputs,
384 .outputs = deint_cuda_outputs,
385 .flags = AVFILTER_FLAG_SUPPORT_TIMELINE_INTERNAL,
386 .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,