From patchwork Fri Nov 2 03:08:38 2018 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Philip Langdale X-Patchwork-Id: 10886 Return-Path: X-Original-To: patchwork@ffaux-bg.ffmpeg.org Delivered-To: patchwork@ffaux-bg.ffmpeg.org Received: from ffbox0-bg.mplayerhq.hu (ffbox0-bg.ffmpeg.org [79.124.17.100]) by ffaux.localdomain (Postfix) with ESMTP id D7BD244DB3F for ; Fri, 2 Nov 2018 05:08:59 +0200 (EET) Received: from [127.0.1.1] (localhost [127.0.0.1]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTP id 5218D68A983; Fri, 2 Nov 2018 05:08:31 +0200 (EET) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org Received: from mail-ua1-f98.google.com (mail-ua1-f98.google.com [209.85.222.98]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id 1459168A948 for ; Fri, 2 Nov 2018 05:08:22 +0200 (EET) Received: by mail-ua1-f98.google.com with SMTP id i30so184734uae.11 for ; Thu, 01 Nov 2018 20:08:54 -0700 (PDT) X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20161025; h=x-gm-message-state:dkim-signature:from:to:cc:subject:date :message-id:in-reply-to:references:mime-version :content-transfer-encoding; bh=VcOHVtDTtAWk0sLBcPX33PcgdYakiS04zshkkRcaZmc=; b=QAAlge3FVt10vAIelY8oPNFvbvDa74TdBEA5lF+b7BPX8aIeqfdXFyr4/ftgKEchdD z83EWOZ47i8Axfod5aHTADeQbZj1b/G0Ng1GqXDneBR5LxtbWXDWa0LkUn2UspwILXsL ZUxbIWb96xWAb8KKPsYLYmWoAUEmL+52i/EbozE/jCHr65ksc0O9xCiJXyxTe1HEISCZ UGIiTZBaAq5CT9OIEMnHywopeBJwfx2QVpXNhds9+7g49X+dxizhs27II2uhKfxXlmNg 8cp78QV5HrvSlVkoWZtfYRzIK4+OXrSO0pQ2TSMlLRRYC1N4jbJmDYCWJlyMbLRVEJaA Dasg== X-Gm-Message-State: AGRZ1gLib8x3yTlwvpfxaeoZb3uzA8yhm0iPe/47D3/CZm1V/X4Uxzzm URDr5FMBtRW0U7NfQ+TkSYwRT2bbhfxECQt0xrEDPx25XausVg== X-Google-Smtp-Source: AJdET5fXCVWOLfyLS+XsdtExqopC6ruQXrjFobuJUMl8fZEK5EtbWTngamaiCprXRKG57UGmzB1VapewgRL7 X-Received: by 2002:a9f:224a:: with SMTP id 68mr4800249uad.38.1541128133426; Thu, 01 Nov 2018 20:08:53 -0700 (PDT) Received: from mail.overt.org (155.208.178.107.bc.googleusercontent.com. [107.178.208.155]) by smtp-relay.gmail.com with ESMTPS id h24sm957381vsj.3.2018.11.01.20.08.53 for (version=TLS1_2 cipher=ECDHE-RSA-AES128-GCM-SHA256 bits=128/128); Thu, 01 Nov 2018 20:08:53 -0700 (PDT) X-Relaying-Domain: gapps.overt.org Received: from authenticated-user (mail.overt.org [107.178.208.155]) (using TLSv1.2 with cipher ECDHE-RSA-AES128-GCM-SHA256 (128/128 bits)) (No client certificate requested) by mail.overt.org (Postfix) with ESMTPSA id 13F5460426; Fri, 2 Nov 2018 03:08:52 +0000 (UTC) DKIM-Signature: v=1; a=rsa-sha256; c=simple/simple; d=overt.org; s=mail; t=1541128132; bh=2nuA9D8ujk2Uj6YCVu0vAms2RdweLkx2lm1Z1wxWWPs=; h=From:To:Cc:Subject:Date:In-Reply-To:References:From; b=onUi7ePz8LlCXQlD5X50l1W7yeBQxiMAcyDYpODHxDaXVpDUitA3ifK1ss0bq2XAA i8nKUmaJxb1bluQebk55NynGJDkx7CMjsPgJ0BjuOI37tcyzsuVWxjuH4+Mqzisv/S +VPWBvRQrgGgtI6mjLbHd3Oe4NP/BUTFZ2PV1QUqD/cAEZMEMw9vbRD4GB0TFHPafe vDswbAhS31Ghjc0h5pYjq+OxK9vURXg1PnAf2BmB7h8WEuY4UAFnImpt57zWbYvyV9 YcOxM2SeKMpGTyE4PV4/4Kzt1QsMbf3y4jPlgOdk4de8XKwPCTWBitiqxWkbikcYHV V4DgErOt/OVzA== From: Philip Langdale To: ffmpeg-devel@ffmpeg.org Date: Thu, 1 Nov 2018 20:08:38 -0700 Message-Id: <20181102030839.8674-3-philipl@overt.org> In-Reply-To: <20181102030839.8674-1-philipl@overt.org> References: <20181102030839.8674-1-philipl@overt.org> MIME-Version: 1.0 Subject: [FFmpeg-devel] [PATCH 2/3] avfilter/vf_yadif_cuda: CUDA accelerated deinterlacer X-BeenThere: ffmpeg-devel@ffmpeg.org X-Mailman-Version: 2.1.20 Precedence: list List-Id: FFmpeg development discussions and patches List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Reply-To: FFmpeg development discussions and patches Cc: Philip Langdale Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" Signed-off-by: Philip Langdale --- Changelog | 1 + configure | 1 + doc/filters.texi | 58 +++++ libavfilter/Makefile | 1 + libavfilter/allfilters.c | 1 + libavfilter/vf_yadif_cuda.c | 426 +++++++++++++++++++++++++++++++++++ libavfilter/vf_yadif_cuda.cu | 296 ++++++++++++++++++++++++ 7 files changed, 784 insertions(+) create mode 100644 libavfilter/vf_yadif_cuda.c create mode 100644 libavfilter/vf_yadif_cuda.cu diff --git a/Changelog b/Changelog index 8430da3c6a..f92886fc2e 100644 --- a/Changelog +++ b/Changelog @@ -44,6 +44,7 @@ version 4.1: - xstack filter - pcm vidc decoder and encoder - (a)graphmonitor filter +- yadif_cuda filter version 4.0: diff --git a/configure b/configure index 2606b885b0..f3fa0cde86 100755 --- a/configure +++ b/configure @@ -3482,6 +3482,7 @@ zscale_filter_deps="libzimg const_nan" scale_vaapi_filter_deps="vaapi" vpp_qsv_filter_deps="libmfx" vpp_qsv_filter_select="qsvvpp" +yadif_cuda_filter_deps="cuda_sdk" # examples avio_dir_cmd_deps="avformat avutil" diff --git a/doc/filters.texi b/doc/filters.texi index 4345a4931b..5d4bfd2e8e 100644 --- a/doc/filters.texi +++ b/doc/filters.texi @@ -17943,6 +17943,64 @@ filter"). It accepts the following parameters: +@table @option + +@item mode +The interlacing mode to adopt. It accepts one of the following values: + +@table @option +@item 0, send_frame +Output one frame for each frame. +@item 1, send_field +Output one frame for each field. +@item 2, send_frame_nospatial +Like @code{send_frame}, but it skips the spatial interlacing check. +@item 3, send_field_nospatial +Like @code{send_field}, but it skips the spatial interlacing check. +@end table + +The default value is @code{send_frame}. + +@item parity +The picture field parity assumed for the input interlaced video. It accepts one +of the following values: + +@table @option +@item 0, tff +Assume the top field is first. +@item 1, bff +Assume the bottom field is first. +@item -1, auto +Enable automatic detection of field parity. +@end table + +The default value is @code{auto}. +If the interlacing is unknown or the decoder does not export this information, +top field first will be assumed. + +@item deint +Specify which frames to deinterlace. Accept one of the following +values: + +@table @option +@item 0, all +Deinterlace all frames. +@item 1, interlaced +Only deinterlace frames marked as interlaced. +@end table + +The default value is @code{all}. +@end table + +@section yadif_cuda + +Deinterlace the input video using the @ref{yadif} algorithm, but implemented +in CUDA so that it can work as part of a GPU accelerated pipeline with nvdec +and/or nvenc. + +It accepts the following parameters: + + @table @option @item mode diff --git a/libavfilter/Makefile b/libavfilter/Makefile index ffbcb40806..4b78b29fad 100644 --- a/libavfilter/Makefile +++ b/libavfilter/Makefile @@ -408,6 +408,7 @@ OBJS-$(CONFIG_WEAVE_FILTER) += vf_weave.o OBJS-$(CONFIG_XBR_FILTER) += vf_xbr.o OBJS-$(CONFIG_XSTACK_FILTER) += vf_stack.o framesync.o OBJS-$(CONFIG_YADIF_FILTER) += vf_yadif.o yadif_common.o +OBJS-$(CONFIG_YADIF_CUDA_FILTER) += vf_yadif_cuda.o vf_yadif_cuda.ptx.o yadif_common.o OBJS-$(CONFIG_ZMQ_FILTER) += f_zmq.o OBJS-$(CONFIG_ZOOMPAN_FILTER) += vf_zoompan.o OBJS-$(CONFIG_ZSCALE_FILTER) += vf_zscale.o diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c index d5a211bda5..c40c7e3a3c 100644 --- a/libavfilter/allfilters.c +++ b/libavfilter/allfilters.c @@ -389,6 +389,7 @@ extern AVFilter ff_vf_weave; extern AVFilter ff_vf_xbr; extern AVFilter ff_vf_xstack; extern AVFilter ff_vf_yadif; +extern AVFilter ff_vf_yadif_cuda; extern AVFilter ff_vf_zmq; extern AVFilter ff_vf_zoompan; extern AVFilter ff_vf_zscale; diff --git a/libavfilter/vf_yadif_cuda.c b/libavfilter/vf_yadif_cuda.c new file mode 100644 index 0000000000..be22344d9d --- /dev/null +++ b/libavfilter/vf_yadif_cuda.c @@ -0,0 +1,426 @@ +/* + * Copyright (C) 2018 Philip Langdale + * + * This file is part of FFmpeg. + * + * FFmpeg is free software; you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public + * License as published by the Free Software Foundation; either + * version 2.1 of the License, or (at your option) any later version. + * + * FFmpeg is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with FFmpeg; if not, write to the Free Software + * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA + */ + +#include +#include "libavutil/avassert.h" +#include "libavutil/hwcontext_cuda.h" +#include "internal.h" +#include "yadif.h" + +extern char vf_yadif_cuda_ptx[]; + +typedef struct DeintCUDAContext { + YADIFContext yadif; + + AVCUDADeviceContext *hwctx; + AVBufferRef *device_ref; + AVBufferRef *input_frames_ref; + AVHWFramesContext *input_frames; + + CUcontext cu_ctx; + CUstream stream; + CUmodule cu_module; + CUfunction cu_func_uchar; + CUfunction cu_func_uchar2; + CUfunction cu_func_ushort; + CUfunction cu_func_ushort2; +} DeintCUDAContext; + +#define DIV_UP(a, b) ( ((a) + (b) - 1) / (b) ) +#define ALIGN_UP(a, b) (((a) + (b) - 1) & ~((b) - 1)) +#define BLOCKX 32 +#define BLOCKY 16 + +static int check_cu(AVFilterContext *avctx, CUresult err, const char *func) +{ + const char *err_name; + const char *err_string; + + av_log(avctx, AV_LOG_TRACE, "Calling %s\n", func); + + if (err == CUDA_SUCCESS) + return 0; + + cuGetErrorName(err, &err_name); + cuGetErrorString(err, &err_string); + + av_log(avctx, AV_LOG_ERROR, "%s failed", func); + if (err_name && err_string) + av_log(avctx, AV_LOG_ERROR, " -> %s: %s", err_name, err_string); + av_log(avctx, AV_LOG_ERROR, "\n"); + + return AVERROR_EXTERNAL; +} + +#define CHECK_CU(x) check_cu(ctx, (x), #x) + +static CUresult call_kernel(AVFilterContext *ctx, CUfunction func, + CUdeviceptr prev, CUdeviceptr cur, CUdeviceptr next, + CUarray_format format, int channels, + int src_width, // Width is pixels per channel + int src_height, // Height is pixels per channel + int src_pitch, // Pitch is bytes + CUdeviceptr dst, + int dst_width, // Width is pixels per channel + int dst_height, // Height is pixels per channel + int dst_pitch, // Pitch is pixels per channel + int parity, int tff) +{ + DeintCUDAContext *s = ctx->priv; + CUtexObject tex_prev = 0, tex_cur = 0, tex_next = 0; + CUresult err; + int skip_spatial_check = s->yadif.mode&2; + + void *args[] = { &dst, &tex_prev, &tex_cur, &tex_next, + &dst_width, &dst_height, &dst_pitch, + &src_width, &src_height, &parity, &tff, + &skip_spatial_check }; + + CUDA_TEXTURE_DESC tex_desc = { + .filterMode = CU_TR_FILTER_MODE_POINT, + .flags = CU_TRSF_READ_AS_INTEGER, + }; + + CUDA_RESOURCE_DESC res_desc = { + .resType = CU_RESOURCE_TYPE_PITCH2D, + .res.pitch2D.format = format, + .res.pitch2D.numChannels = channels, + .res.pitch2D.width = src_width, + .res.pitch2D.height = src_height, + .res.pitch2D.pitchInBytes = src_pitch, + }; + + res_desc.res.pitch2D.devPtr = (CUdeviceptr)prev; + err = CHECK_CU(cuTexObjectCreate(&tex_prev, &res_desc, &tex_desc, NULL)); + if (err != CUDA_SUCCESS) { + goto exit; + } + + res_desc.res.pitch2D.devPtr = (CUdeviceptr)cur; + err = CHECK_CU(cuTexObjectCreate(&tex_cur, &res_desc, &tex_desc, NULL)); + if (err != CUDA_SUCCESS) { + goto exit; + } + + res_desc.res.pitch2D.devPtr = (CUdeviceptr)next; + err = CHECK_CU(cuTexObjectCreate(&tex_next, &res_desc, &tex_desc, NULL)); + if (err != CUDA_SUCCESS) { + goto exit; + } + + err = CHECK_CU(cuLaunchKernel(func, + DIV_UP(dst_width, BLOCKX), DIV_UP(dst_height, BLOCKY), 1, + BLOCKX, BLOCKY, 1, + 0, s->stream, args, NULL)); + +exit: + if (tex_prev) + CHECK_CU(cuTexObjectDestroy(tex_prev)); + if (tex_cur) + CHECK_CU(cuTexObjectDestroy(tex_cur)); + if (tex_next) + CHECK_CU(cuTexObjectDestroy(tex_next)); + + return err; +} + +static void filter(AVFilterContext *ctx, AVFrame *dst, + int parity, int tff) +{ + DeintCUDAContext *s = ctx->priv; + YADIFContext *y = &s->yadif; + CUcontext dummy; + CUresult err; + int i; + + err = CHECK_CU(cuCtxPushCurrent(s->cu_ctx)); + if (err != CUDA_SUCCESS) { + goto exit; + } + + for (i = 0; i < y->csp->nb_components; i++) { + CUfunction func; + CUarray_format format; + int pixel_size, channels; + const AVComponentDescriptor *comp = &y->csp->comp[i]; + + if (comp->plane < i) { + // We process planes as a whole, so don't reprocess + // them for additional components + continue; + } + + pixel_size = (comp->depth + comp->shift) / 8; + channels = comp->step / pixel_size; + if (pixel_size > 2 || channels > 2) { + av_log(ctx, AV_LOG_ERROR, "Unsupported pixel format: %s\n", y->csp->name); + goto exit; + } + switch (pixel_size) { + case 1: + func = channels == 1 ? s->cu_func_uchar : s->cu_func_uchar2; + format = CU_AD_FORMAT_UNSIGNED_INT8; + break; + case 2: + func = channels == 1 ? s->cu_func_ushort : s->cu_func_ushort2; + format = CU_AD_FORMAT_UNSIGNED_INT16; + break; + default: + av_log(ctx, AV_LOG_ERROR, "Unsupported pixel format: %s\n", y->csp->name); + goto exit; + } + av_log(ctx, AV_LOG_TRACE, + "Deinterlacing plane %d: pixel_size: %d channels: %d\n", + comp->plane, pixel_size, channels); + call_kernel(ctx, func, + (CUdeviceptr)y->prev->data[i], + (CUdeviceptr)y->cur->data[i], + (CUdeviceptr)y->next->data[i], + format, channels, + AV_CEIL_RSHIFT(y->cur->width, i ? y->csp->log2_chroma_w : 0), + AV_CEIL_RSHIFT(y->cur->height, i ? y->csp->log2_chroma_h : 0), + y->cur->linesize[i], + (CUdeviceptr)dst->data[i], + AV_CEIL_RSHIFT(dst->width, i ? y->csp->log2_chroma_w : 0), + AV_CEIL_RSHIFT(dst->height, i ? y->csp->log2_chroma_h : 0), + dst->linesize[i] / comp->step, + parity, tff); + } + + err = CHECK_CU(cuStreamSynchronize(s->stream)); + if (err != CUDA_SUCCESS) { + goto exit; + } + +exit: + CHECK_CU(cuCtxPopCurrent(&dummy)); + return; +} + +static av_cold void deint_cuda_uninit(AVFilterContext *ctx) +{ + CUcontext dummy; + DeintCUDAContext *s = ctx->priv; + YADIFContext *y = &s->yadif; + + if (s->cu_module) { + CHECK_CU(cuCtxPushCurrent(s->cu_ctx)); + CHECK_CU(cuModuleUnload(s->cu_module)); + CHECK_CU(cuCtxPopCurrent(&dummy)); + } + + av_frame_free(&y->prev); + av_frame_free(&y->cur); + av_frame_free(&y->next); + + av_buffer_unref(&s->device_ref); + s->hwctx = NULL; + av_buffer_unref(&s->input_frames_ref); + s->input_frames = NULL; +} + +static int deint_cuda_query_formats(AVFilterContext *ctx) +{ + enum AVPixelFormat pix_fmts[] = { + AV_PIX_FMT_CUDA, AV_PIX_FMT_NONE, + }; + int ret; + + if ((ret = ff_formats_ref(ff_make_format_list(pix_fmts), + &ctx->inputs[0]->out_formats)) < 0) + return ret; + if ((ret = ff_formats_ref(ff_make_format_list(pix_fmts), + &ctx->outputs[0]->in_formats)) < 0) + return ret; + + return 0; +} + +static int config_input(AVFilterLink *inlink) +{ + AVFilterContext *ctx = inlink->dst; + DeintCUDAContext *s = ctx->priv; + + if (!inlink->hw_frames_ctx) { + av_log(ctx, AV_LOG_ERROR, "A hardware frames reference is " + "required to associate the processing device.\n"); + return AVERROR(EINVAL); + } + + s->input_frames_ref = av_buffer_ref(inlink->hw_frames_ctx); + if (!s->input_frames_ref) { + av_log(ctx, AV_LOG_ERROR, "A input frames reference create " + "failed.\n"); + return AVERROR(ENOMEM); + } + s->input_frames = (AVHWFramesContext*)s->input_frames_ref->data; + + return 0; +} + +static int config_output(AVFilterLink *link) +{ + AVHWFramesContext *output_frames; + AVFilterContext *ctx = link->src; + DeintCUDAContext *s = ctx->priv; + YADIFContext *y = &s->yadif; + int ret = 0; + CUcontext dummy; + CUresult err; + + av_assert0(s->input_frames); + s->device_ref = av_buffer_ref(s->input_frames->device_ref); + if (!s->device_ref) { + av_log(ctx, AV_LOG_ERROR, "A device reference create " + "failed.\n"); + return AVERROR(ENOMEM); + } + s->hwctx = ((AVHWDeviceContext*)s->device_ref->data)->hwctx; + s->cu_ctx = s->hwctx->cuda_ctx; + s->stream = s->hwctx->stream; + + link->hw_frames_ctx = av_hwframe_ctx_alloc(s->device_ref); + if (!link->hw_frames_ctx) { + av_log(ctx, AV_LOG_ERROR, "Failed to create HW frame context " + "for output.\n"); + ret = AVERROR(ENOMEM); + goto exit; + } + + output_frames = (AVHWFramesContext*)link->hw_frames_ctx->data; + + output_frames->format = AV_PIX_FMT_CUDA; + output_frames->sw_format = s->input_frames->sw_format; + output_frames->width = ctx->inputs[0]->w; + output_frames->height = ctx->inputs[0]->h; + + output_frames->initial_pool_size = 4; + + ret = ff_filter_init_hw_frames(ctx, link, 10); + if (ret < 0) + goto exit; + + ret = av_hwframe_ctx_init(link->hw_frames_ctx); + if (ret < 0) { + av_log(ctx, AV_LOG_ERROR, "Failed to initialise CUDA frame " + "context for output: %d\n", ret); + goto exit; + } + + link->time_base.num = ctx->inputs[0]->time_base.num; + link->time_base.den = ctx->inputs[0]->time_base.den * 2; + link->w = ctx->inputs[0]->w; + link->h = ctx->inputs[0]->h; + + if(y->mode & 1) + link->frame_rate = av_mul_q(ctx->inputs[0]->frame_rate, + (AVRational){2, 1}); + + if (link->w < 3 || link->h < 3) { + av_log(ctx, AV_LOG_ERROR, "Video of less than 3 columns or lines is not supported\n"); + ret = AVERROR(EINVAL); + goto exit; + } + + y->csp = av_pix_fmt_desc_get(output_frames->sw_format); + y->filter = filter; + + err = CHECK_CU(cuCtxPushCurrent(s->cu_ctx)); + if (err != CUDA_SUCCESS) { + ret = AVERROR_EXTERNAL; + goto exit; + } + + err = CHECK_CU(cuModuleLoadData(&s->cu_module, vf_yadif_cuda_ptx)); + if (err != CUDA_SUCCESS) { + ret = AVERROR_INVALIDDATA; + goto exit; + } + + err = CHECK_CU(cuModuleGetFunction(&s->cu_func_uchar, s->cu_module, "yadif_uchar")); + if (err != CUDA_SUCCESS) { + ret = AVERROR_INVALIDDATA; + goto exit; + } + + err = CHECK_CU(cuModuleGetFunction(&s->cu_func_uchar2, s->cu_module, "yadif_uchar2")); + if (err != CUDA_SUCCESS) { + ret = AVERROR_INVALIDDATA; + goto exit; + } + + err= CHECK_CU(cuModuleGetFunction(&s->cu_func_ushort, s->cu_module, "yadif_ushort")); + if (err != CUDA_SUCCESS) { + ret = AVERROR_INVALIDDATA; + goto exit; + } + + err = CHECK_CU(cuModuleGetFunction(&s->cu_func_ushort2, s->cu_module, "yadif_ushort2")); + if (err != CUDA_SUCCESS) { + ret = AVERROR_INVALIDDATA; + goto exit; + } + +exit: + CHECK_CU(cuCtxPopCurrent(&dummy)); + + return ret; +} + +static const AVClass yadif_cuda_class = { + .class_name = "yadif_cuda", + .item_name = av_default_item_name, + .option = ff_yadif_options, + .version = LIBAVUTIL_VERSION_INT, + .category = AV_CLASS_CATEGORY_FILTER, +}; + +static const AVFilterPad deint_cuda_inputs[] = { + { + .name = "default", + .type = AVMEDIA_TYPE_VIDEO, + .filter_frame = ff_yadif_filter_frame, + .config_props = config_input, + }, + { NULL } +}; + +static const AVFilterPad deint_cuda_outputs[] = { + { + .name = "default", + .type = AVMEDIA_TYPE_VIDEO, + .request_frame = ff_yadif_request_frame, + .config_props = config_output, + }, + { NULL } +}; + +AVFilter ff_vf_yadif_cuda = { + .name = "yadif_cuda", + .description = NULL_IF_CONFIG_SMALL("Deinterlace CUDA frames"), + .priv_size = sizeof(DeintCUDAContext), + .priv_class = &yadif_cuda_class, + .uninit = deint_cuda_uninit, + .query_formats = deint_cuda_query_formats, + .inputs = deint_cuda_inputs, + .outputs = deint_cuda_outputs, + .flags = AVFILTER_FLAG_SUPPORT_TIMELINE_INTERNAL, + .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE, +}; diff --git a/libavfilter/vf_yadif_cuda.cu b/libavfilter/vf_yadif_cuda.cu new file mode 100644 index 0000000000..65a902c66b --- /dev/null +++ b/libavfilter/vf_yadif_cuda.cu @@ -0,0 +1,296 @@ +/* + * Copyright (C) 2018 Philip Langdale + * + * This file is part of FFmpeg. + * + * FFmpeg is free software; you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public + * License as published by the Free Software Foundation; either + * version 2.1 of the License, or (at your option) any later version. + * + * FFmpeg is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with FFmpeg; if not, write to the Free Software + * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA + */ + +template +__inline__ __device__ T spatial_predictor(T a, T b, T c, T d, T e, T f, T g, + T h, T i, T j, T k, T l, T m, T n) +{ + int spatial_pred = (d + k)/2; + int spatial_score = abs(c - j) + abs(d - k) + abs(e - l); + + int score = abs(b - k) + abs(c - l) + abs(d - m); + if (score < spatial_score) { + spatial_pred = (c + l)/2; + spatial_score = score; + score = abs(a - l) + abs(b - m) + abs(c - n); + if (score < spatial_score) { + spatial_pred = (b + m)/2; + spatial_score = score; + } + } + score = abs(d - i) + abs(e - j) + abs(f - k); + if (score < spatial_score) { + spatial_pred = (e + j)/2; + spatial_score = score; + score = abs(e - h) + abs(f - i) + abs(g - j); + if (score < spatial_score) { + spatial_pred = (f + i)/2; + spatial_score = score; + } + } + return spatial_pred; +} + +__inline__ __device__ int max3(int a, int b, int c) +{ + int x = max(a, b); + return max(x, c); +} + +__inline__ __device__ int min3(int a, int b, int c) +{ + int x = min(a, b); + return min(x, c); +} + +template +__inline__ __device__ T temporal_predictor(T A, T B, T C, T D, T E, T F, + T G, T H, T I, T J, T K, T L, + T spatial_pred, bool skip_check) +{ + int p0 = (C + H) / 2; + int p1 = F; + int p2 = (D + I) / 2; + int p3 = G; + int p4 = (E + J) / 2; + + int tdiff0 = abs(D - I); + int tdiff1 = (abs(A - F) + abs(B - G)) / 2; + int tdiff2 = (abs(K - F) + abs(G - L)) / 2; + + int diff = max3(tdiff0, tdiff1, tdiff2); + + if (!skip_check) { + int maxi = max3(p2 - p3, p2 - p1, min(p0 - p1, p4 - p3)); + int mini = min3(p2 - p3, p2 - p1, max(p0 - p1, p4 - p3)); + diff = max3(diff, mini, -maxi); + } + + if (spatial_pred > p2 + diff) { + spatial_pred = p2 + diff; + } + if (spatial_pred < p2 - diff) { + spatial_pred = p2 - diff; + } + + return spatial_pred; +} + +template +__inline__ __device__ void yadif_single(T *dst, + cudaTextureObject_t prev, + cudaTextureObject_t cur, + cudaTextureObject_t next, + int dst_width, int dst_height, int dst_pitch, + int src_width, int src_height, + int parity, int tff, bool skip_spatial_check) +{ + // Identify location + int xo = blockIdx.x * blockDim.x + threadIdx.x; + int yo = blockIdx.y * blockDim.y + threadIdx.y; + + if (xo >= dst_width || yo >= dst_height) { + return; + } + + // Don't modify the primary field + if (yo % 2 == parity) { + dst[yo*dst_pitch+xo] = tex2D(cur, xo, yo); + return; + } + + // Calculate spatial prediction + T a = tex2D(cur, xo - 3, yo - 1); + T b = tex2D(cur, xo - 2, yo - 1); + T c = tex2D(cur, xo - 1, yo - 1); + T d = tex2D(cur, xo - 0, yo - 1); + T e = tex2D(cur, xo + 1, yo - 1); + T f = tex2D(cur, xo + 2, yo - 1); + T g = tex2D(cur, xo + 3, yo - 1); + + T h = tex2D(cur, xo - 3, yo + 1); + T i = tex2D(cur, xo - 2, yo + 1); + T j = tex2D(cur, xo - 1, yo + 1); + T k = tex2D(cur, xo - 0, yo + 1); + T l = tex2D(cur, xo + 1, yo + 1); + T m = tex2D(cur, xo + 2, yo + 1); + T n = tex2D(cur, xo + 3, yo + 1); + + T spatial_pred = + spatial_predictor(a, b, c, d, e, f, g, h, i, j, k, l, m, n); + + // Calculate temporal prediction + int is_second_field = !(parity ^ tff); + + cudaTextureObject_t prev2 = prev; + cudaTextureObject_t prev1 = is_second_field ? cur : prev; + cudaTextureObject_t next1 = is_second_field ? next : cur; + cudaTextureObject_t next2 = next; + + T A = tex2D(prev2, xo, yo - 1); + T B = tex2D(prev2, xo, yo + 1); + T C = tex2D(prev1, xo, yo - 2); + T D = tex2D(prev1, xo, yo + 0); + T E = tex2D(prev1, xo, yo + 2); + T F = tex2D(cur, xo, yo - 1); + T G = tex2D(cur, xo, yo + 1); + T H = tex2D(next1, xo, yo - 2); + T I = tex2D(next1, xo, yo + 0); + T J = tex2D(next1, xo, yo + 2); + T K = tex2D(next2, xo, yo - 1); + T L = tex2D(next2, xo, yo + 1); + + spatial_pred = temporal_predictor(A, B, C, D, E, F, G, H, I, J, K, L, + spatial_pred, skip_spatial_check); + + dst[yo*dst_pitch+xo] = spatial_pred; +} + +template +__inline__ __device__ void yadif_double(T *dst, + cudaTextureObject_t prev, + cudaTextureObject_t cur, + cudaTextureObject_t next, + int dst_width, int dst_height, int dst_pitch, + int src_width, int src_height, + int parity, int tff, bool skip_spatial_check) +{ + int xo = blockIdx.x * blockDim.x + threadIdx.x; + int yo = blockIdx.y * blockDim.y + threadIdx.y; + + if (xo >= dst_width || yo >= dst_height) { + return; + } + + if (yo % 2 == parity) { + // Don't modify the primary field + dst[yo*dst_pitch+xo] = tex2D(cur, xo, yo); + return; + } + + T a = tex2D(cur, xo - 3, yo - 1); + T b = tex2D(cur, xo - 2, yo - 1); + T c = tex2D(cur, xo - 1, yo - 1); + T d = tex2D(cur, xo - 0, yo - 1); + T e = tex2D(cur, xo + 1, yo - 1); + T f = tex2D(cur, xo + 2, yo - 1); + T g = tex2D(cur, xo + 3, yo - 1); + + T h = tex2D(cur, xo - 3, yo + 1); + T i = tex2D(cur, xo - 2, yo + 1); + T j = tex2D(cur, xo - 1, yo + 1); + T k = tex2D(cur, xo - 0, yo + 1); + T l = tex2D(cur, xo + 1, yo + 1); + T m = tex2D(cur, xo + 2, yo + 1); + T n = tex2D(cur, xo + 3, yo + 1); + + T spatial_pred = { + 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), + 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) }; + + // Calculate temporal prediction + int is_second_field = !(parity ^ tff); + + cudaTextureObject_t prev2 = prev; + cudaTextureObject_t prev1 = is_second_field ? cur : prev; + cudaTextureObject_t next1 = is_second_field ? next : cur; + cudaTextureObject_t next2 = next; + + T A = tex2D(prev2, xo, yo - 1); + T B = tex2D(prev2, xo, yo + 1); + T C = tex2D(prev1, xo, yo - 2); + T D = tex2D(prev1, xo, yo + 0); + T E = tex2D(prev1, xo, yo + 2); + T F = tex2D(cur, xo, yo - 1); + T G = tex2D(cur, xo, yo + 1); + T H = tex2D(next1, xo, yo - 2); + T I = tex2D(next1, xo, yo + 0); + T J = tex2D(next1, xo, yo + 2); + T K = tex2D(next2, xo, yo - 1); + T L = tex2D(next2, xo, yo + 1); + + spatial_pred = { + 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, + spatial_pred.x, skip_spatial_check), + 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, + spatial_pred.y, skip_spatial_check) }; + + dst[yo*dst_pitch+xo] = spatial_pred; +} + +extern "C" { + +__global__ void yadif_uchar(unsigned char *dst, + cudaTextureObject_t prev, + cudaTextureObject_t cur, + cudaTextureObject_t next, + int dst_width, int dst_height, int dst_pitch, + int src_width, int src_height, + int parity, int tff, bool skip_spatial_check) +{ + yadif_single(dst, prev, cur, next, + dst_width, dst_height, dst_pitch, + src_width, src_height, + parity, tff, skip_spatial_check); +} + +__global__ void yadif_ushort(unsigned short *dst, + cudaTextureObject_t prev, + cudaTextureObject_t cur, + cudaTextureObject_t next, + int dst_width, int dst_height, int dst_pitch, + int src_width, int src_height, + int parity, int tff, bool skip_spatial_check) +{ + yadif_single(dst, prev, cur, next, + dst_width, dst_height, dst_pitch, + src_width, src_height, + parity, tff, skip_spatial_check); +} + +__global__ void yadif_uchar2(uchar2 *dst, + cudaTextureObject_t prev, + cudaTextureObject_t cur, + cudaTextureObject_t next, + int dst_width, int dst_height, int dst_pitch, + int src_width, int src_height, + int parity, int tff, bool skip_spatial_check) +{ + yadif_double(dst, prev, cur, next, + dst_width, dst_height, dst_pitch, + src_width, src_height, + parity, tff, skip_spatial_check); +} + +__global__ void yadif_ushort2(ushort2 *dst, + cudaTextureObject_t prev, + cudaTextureObject_t cur, + cudaTextureObject_t next, + int dst_width, int dst_height, int dst_pitch, + int src_width, int src_height, + int parity, int tff, bool skip_spatial_check) +{ + yadif_double(dst, prev, cur, next, + dst_width, dst_height, dst_pitch, + src_width, src_height, + parity, tff, skip_spatial_check); +} + +} /* extern "C" */