From patchwork Fri May 12 19:40:16 2017 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Timo Rothenpieler X-Patchwork-Id: 3664 Delivered-To: ffmpegpatchwork@gmail.com Received: by 10.103.3.129 with SMTP id 123csp474397vsd; Fri, 12 May 2017 12:40:51 -0700 (PDT) X-Received: by 10.28.17.21 with SMTP id 21mr4334312wmr.83.1494618051215; Fri, 12 May 2017 12:40:51 -0700 (PDT) ARC-Seal: i=1; a=rsa-sha256; t=1494618051; cv=none; d=google.com; s=arc-20160816; b=mngqqWw+ofwYhEBgUkFFjUhrsUA0ukINt9lNS1ERO1YaHqxuGPDI1PCk0Zx8b+PSfk KIrPsXjs5CMAmxdYzlLovrRCK7Pn2f3CGaKV149qhIo+kR/lXozpb2nMiKIPLtFXRzIA MdBKWSKY7w8iZ0hKixUY6U6GQSKZFnbJ2AYAr3VBh/uxDOedX80sRixwLGMiseQTU+2T Ml6zbxU62gA/sp7i+n95qod6VtbGruen4ZIBs7UfM1OzsQ/n6uY8qo0Taw4w5xpa9aEZ 06NPqYhjLyDPSwdcvJnXoq6VJ0CU6wM8lONbOBb0EriaYMEzBsdqVWQ/ziSFfmIVPUKd x1bA== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=sender:errors-to:content-transfer-encoding:mime-version:cc:reply-to :list-subscribe:list-help:list-post:list-archive:list-unsubscribe :list-id:precedence:subject:references:in-reply-to:message-id:date :to:from:dkim-signature:delivered-to:arc-authentication-results; bh=r0Z3v/0xIkkob8vbqZHj51xqxNHlFCt+X4ObRsRXtn4=; b=TY5l2oOhE3khcsql1/wKZMTQY798bZs/JcjiJfk/T0avmbgUz+mRx46C062bmvA6rn eLJIpYXTXOCNBCHk4h15mk/PtM8XN7BhkqOnUoEowG3kUFl92EYXdZjmuEXJiW68rcNB SvWE2YkxwX2sGXWa3dnppKphXwjthdLbpe6meVpylsFMVE98CZvCvKR4MAa23SgHXrVo IB58qX0QryIw7QCJJCN+RCytVeGzqNLANbY+HvleK5jV8KzjPM1Zm0wPhnH87YjUU7eU I4BK4YJ4rtyvTgA1OpcN4PeW+ROuEv7VYbfmNJd2pt7XoXftGaIYA8QVlNiaNer2J7N+ j3ZQ== ARC-Authentication-Results: i=1; mx.google.com; dkim=neutral (body hash did not verify) header.i=@rothenpieler.org; spf=pass (google.com: domain of ffmpeg-devel-bounces@ffmpeg.org designates 79.124.17.100 as permitted sender) smtp.mailfrom=ffmpeg-devel-bounces@ffmpeg.org Return-Path: Received: from ffbox0-bg.mplayerhq.hu (ffbox0-bg.ffmpeg.org. [79.124.17.100]) by mx.google.com with ESMTP id n202si1758859wmg.10.2017.05.12.12.40.50; Fri, 12 May 2017 12:40:51 -0700 (PDT) Received-SPF: pass (google.com: domain of ffmpeg-devel-bounces@ffmpeg.org designates 79.124.17.100 as permitted sender) client-ip=79.124.17.100; Authentication-Results: mx.google.com; dkim=neutral (body hash did not verify) header.i=@rothenpieler.org; spf=pass (google.com: domain of ffmpeg-devel-bounces@ffmpeg.org designates 79.124.17.100 as permitted sender) smtp.mailfrom=ffmpeg-devel-bounces@ffmpeg.org Received: from [127.0.1.1] (localhost [127.0.0.1]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTP id CD78B6883BE; Fri, 12 May 2017 22:40:24 +0300 (EEST) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org Received: from btbn.de (btbn.de [5.9.118.179]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id 8C5BE680418 for ; Fri, 12 May 2017 22:40:17 +0300 (EEST) Received: from localhost.localdomain (unknown [IPv6:2a02:8109:43f:959c:ba97:5aff:fe10:ec69]) by btbn.de (Postfix) with ESMTPSA id 089516C416; Fri, 12 May 2017 21:40:25 +0200 (CEST) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/simple; d=rothenpieler.org; s=mail; t=1494618025; bh=7CGJo8w5M+kbJZFRabNkG+h1dbymDH4nFHwwGPOHRPI=; h=From:To:Cc:Subject:Date:In-Reply-To:References; b=kFv/VbY54EP9oCUqnHlJADsQRCGdgcU5HG0xFY+gXGdbKPJWLGTHFSvomovyaSqfh YRjXOZaM9YBWbFv7zqRYZib5WI0ymQXueceSEstj3NGF9B4IjSOWt8szOpmBocA5W3 XkdILk7aTv5f1mAs62qoPHqbJuCwTsnVtyhmOJ9k= From: Timo Rothenpieler To: ffmpeg-devel@ffmpeg.org Date: Fri, 12 May 2017 21:40:16 +0200 Message-Id: <20170512194016.8298-3-timo@rothenpieler.org> X-Mailer: git-send-email 2.12.2 In-Reply-To: <20170512194016.8298-1-timo@rothenpieler.org> References: <20170511205921.2172-1-timo@rothenpieler.org> <20170512194016.8298-1-timo@rothenpieler.org> Subject: [FFmpeg-devel] [PATCH 3/3] avfilter/scale_cuda: add CUDA scale filter 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: Yogender Gupta MIME-Version: 1.0 Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" From: Yogender Gupta --- libavfilter/Makefile | 1 + libavfilter/allfilters.c | 1 + libavfilter/vf_scale_cuda.c | 555 +++++++++++++++++++++++++++++++++++++++++++ libavfilter/vf_scale_cuda.cu | 212 +++++++++++++++++ 4 files changed, 769 insertions(+) create mode 100644 libavfilter/vf_scale_cuda.c create mode 100644 libavfilter/vf_scale_cuda.cu diff --git a/libavfilter/Makefile b/libavfilter/Makefile index f7dfe8ad54..f177fdb42b 100644 --- a/libavfilter/Makefile +++ b/libavfilter/Makefile @@ -267,6 +267,7 @@ OBJS-$(CONFIG_REVERSE_FILTER) += f_reverse.o OBJS-$(CONFIG_ROTATE_FILTER) += vf_rotate.o OBJS-$(CONFIG_SAB_FILTER) += vf_sab.o OBJS-$(CONFIG_SCALE_FILTER) += vf_scale.o scale.o +OBJS-$(CONFIG_SCALE_CUDA_FILTER) += vf_scale_cuda.o vf_scale_cuda.ptx.o OBJS-$(CONFIG_SCALE_NPP_FILTER) += vf_scale_npp.o scale.o OBJS-$(CONFIG_SCALE_QSV_FILTER) += vf_scale_qsv.o OBJS-$(CONFIG_SCALE_VAAPI_FILTER) += vf_scale_vaapi.o scale.o diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c index cd35ae4c9c..a8939b9094 100644 --- a/libavfilter/allfilters.c +++ b/libavfilter/allfilters.c @@ -278,6 +278,7 @@ static void register_all(void) REGISTER_FILTER(ROTATE, rotate, vf); REGISTER_FILTER(SAB, sab, vf); REGISTER_FILTER(SCALE, scale, vf); + REGISTER_FILTER(SCALE_CUDA, scale_cuda, vf); REGISTER_FILTER(SCALE_NPP, scale_npp, vf); REGISTER_FILTER(SCALE_QSV, scale_qsv, vf); REGISTER_FILTER(SCALE_VAAPI, scale_vaapi, vf); diff --git a/libavfilter/vf_scale_cuda.c b/libavfilter/vf_scale_cuda.c new file mode 100644 index 0000000000..1f643197ac --- /dev/null +++ b/libavfilter/vf_scale_cuda.c @@ -0,0 +1,555 @@ +/* +* Copyright (c) 2017, NVIDIA CORPORATION. All rights reserved. +* +* Permission is hereby granted, free of charge, to any person obtaining a +* copy of this software and associated documentation files (the "Software"), +* to deal in the Software without restriction, including without limitation +* the rights to use, copy, modify, merge, publish, distribute, sublicense, +* and/or sell copies of the Software, and to permit persons to whom the +* Software is furnished to do so, subject to the following conditions: +* +* The above copyright notice and this permission notice shall be included in +* all copies or substantial portions of the Software. +* +* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL +* THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING +* FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER +* DEALINGS IN THE SOFTWARE. +*/ + +#include +#include +#include + +#include "libavutil/avstring.h" +#include "libavutil/common.h" +#include "libavutil/hwcontext.h" +#include "libavutil/hwcontext_cuda_internal.h" +#include "libavutil/internal.h" +#include "libavutil/opt.h" +#include "libavutil/pixdesc.h" + +#include "avfilter.h" +#include "formats.h" +#include "internal.h" +#include "scale.h" +#include "video.h" + +static const enum AVPixelFormat supported_formats[] = { + AV_PIX_FMT_YUV420P, + AV_PIX_FMT_NV12, + AV_PIX_FMT_YUV444P, + AV_PIX_FMT_P010, + AV_PIX_FMT_P016 +}; + +#define DIV_UP(a, b) ( ((a) + (b) - 1) / (b) ) +#define ALIGN_UP(a, b) ((a + b -1) & ~(b-1)) +#define NUM_BUFFERS 2 +#define BLOCKX 32 +#define BLOCKY 16 + +typedef struct CUDAScaleContext { + const AVClass *class; + enum AVPixelFormat in_fmt; + enum AVPixelFormat out_fmt; + + struct { + int width; + int height; + } planes_in[3], planes_out[3]; + + AVBufferRef *frames_ctx; + AVFrame *frame; + + AVFrame *tmp_frame; + int passthrough; + + /** + * Output sw format. AV_PIX_FMT_NONE for no conversion. + */ + enum AVPixelFormat format; + + char *w_expr; ///< width expression string + char *h_expr; ///< height expression string + + CUcontext cu_ctx; + CUevent cu_event; + CUmodule cu_module; + CUfunction cu_func_uchar; + CUfunction cu_func_uchar2; + CUfunction cu_func_uchar4; + CUfunction cu_func_ushort; + CUfunction cu_func_ushort2; + CUfunction cu_func_ushort4; + CUtexref cu_tex_uchar; + CUtexref cu_tex_uchar2; + CUtexref cu_tex_uchar4; + CUtexref cu_tex_ushort; + CUtexref cu_tex_ushort2; + CUtexref cu_tex_ushort4; + + CUdeviceptr srcBuffer; + CUdeviceptr dstBuffer; + int tex_alignment; +} CUDAScaleContext; + +static int cudascale_init(AVFilterContext *ctx) +{ + CUDAScaleContext *s = ctx->priv; + + s->format = AV_PIX_FMT_NONE; + s->frame = av_frame_alloc(); + if (!s->frame) + return AVERROR(ENOMEM); + + s->tmp_frame = av_frame_alloc(); + if (!s->tmp_frame) + return AVERROR(ENOMEM); + + return 0; +} + +static void cudascale_uninit(AVFilterContext *ctx) +{ + CUDAScaleContext *s = ctx->priv; + + av_frame_free(&s->frame); + av_buffer_unref(&s->frames_ctx); + av_frame_free(&s->tmp_frame); +} + +static int cudascale_query_formats(AVFilterContext *ctx) +{ + static const enum AVPixelFormat pixel_formats[] = { + AV_PIX_FMT_CUDA, AV_PIX_FMT_NONE, + }; + AVFilterFormats *pix_fmts = ff_make_format_list(pixel_formats); + + return ff_set_common_formats(ctx, pix_fmts); +} + +static int init_stage(CUDAScaleContext *s, AVBufferRef *device_ctx) +{ + AVBufferRef *out_ref = NULL; + AVHWFramesContext *out_ctx; + int in_sw, in_sh, out_sw, out_sh; + int ret, i; + + av_pix_fmt_get_chroma_sub_sample(s->in_fmt, &in_sw, &in_sh); + av_pix_fmt_get_chroma_sub_sample(s->out_fmt, &out_sw, &out_sh); + if (!s->planes_out[0].width) { + s->planes_out[0].width = s->planes_in[0].width; + s->planes_out[0].height = s->planes_in[0].height; + } + + for (i = 1; i < FF_ARRAY_ELEMS(s->planes_in); i++) { + s->planes_in[i].width = s->planes_in[0].width >> in_sw; + s->planes_in[i].height = s->planes_in[0].height >> in_sh; + s->planes_out[i].width = s->planes_out[0].width >> out_sw; + s->planes_out[i].height = s->planes_out[0].height >> out_sh; + } + + out_ref = av_hwframe_ctx_alloc(device_ctx); + if (!out_ref) + return AVERROR(ENOMEM); + out_ctx = (AVHWFramesContext*)out_ref->data; + + out_ctx->format = AV_PIX_FMT_CUDA; + out_ctx->sw_format = s->out_fmt; + out_ctx->width = FFALIGN(s->planes_out[0].width, 32); + out_ctx->height = FFALIGN(s->planes_out[0].height, 32); + + ret = av_hwframe_ctx_init(out_ref); + if (ret < 0) + goto fail; + + av_frame_unref(s->frame); + ret = av_hwframe_get_buffer(out_ref, s->frame, 0); + if (ret < 0) + goto fail; + + s->frame->width = s->planes_out[0].width; + s->frame->height = s->planes_out[0].height; + + av_buffer_unref(&s->frames_ctx); + s->frames_ctx = out_ref; + + return 0; +fail: + av_buffer_unref(&out_ref); + return ret; +} + +static int format_is_supported(enum AVPixelFormat fmt) +{ + int i; + + for (i = 0; i < FF_ARRAY_ELEMS(supported_formats); i++) + if (supported_formats[i] == fmt) + return 1; + return 0; +} + +static int init_processing_chain(AVFilterContext *ctx, int in_width, int in_height, + int out_width, int out_height) +{ + CUDAScaleContext *s = ctx->priv; + + AVHWFramesContext *in_frames_ctx; + + enum AVPixelFormat in_format; + enum AVPixelFormat out_format; + int ret; + + /* check that we have a hw context */ + if (!ctx->inputs[0]->hw_frames_ctx) { + av_log(ctx, AV_LOG_ERROR, "No hw context provided on input\n"); + return AVERROR(EINVAL); + } + in_frames_ctx = (AVHWFramesContext*)ctx->inputs[0]->hw_frames_ctx->data; + in_format = in_frames_ctx->sw_format; + out_format = (s->format == AV_PIX_FMT_NONE) ? in_format : s->format; + + if (!format_is_supported(in_format)) { + av_log(ctx, AV_LOG_ERROR, "Unsupported input format: %s\n", + av_get_pix_fmt_name(in_format)); + return AVERROR(ENOSYS); + } + if (!format_is_supported(out_format)) { + av_log(ctx, AV_LOG_ERROR, "Unsupported output format: %s\n", + av_get_pix_fmt_name(out_format)); + return AVERROR(ENOSYS); + } + + if (in_width == out_width && in_height == out_height) + s->passthrough = 1; + + s->in_fmt = in_format; + s->out_fmt = out_format; + + s->planes_in[0].width = in_width; + s->planes_in[0].height = in_height; + s->planes_out[0].width = out_width; + s->planes_out[0].height = out_height; + + ret = init_stage(s, in_frames_ctx->device_ref); + if (ret < 0) + return ret; + + ctx->outputs[0]->hw_frames_ctx = av_buffer_ref(s->frames_ctx); + if (!ctx->outputs[0]->hw_frames_ctx) + return AVERROR(ENOMEM); + + return 0; +} + +static int cudascale_config_props(AVFilterLink *outlink) +{ + AVFilterContext *ctx = outlink->src; + AVFilterLink *inlink = outlink->src->inputs[0]; + CUDAScaleContext *s = ctx->priv; + AVHWFramesContext *frames_ctx = (AVHWFramesContext*)inlink->hw_frames_ctx->data; + AVCUDADeviceContext *device_hwctx = frames_ctx->device_ctx->hwctx; + CUcontext dummy, cuda_ctx = device_hwctx->cuda_ctx; + CUresult err; + int w, h; + int ret; + + extern char vf_scale_cuda_ptx[]; + + err = cuCtxPushCurrent(cuda_ctx); + if (err != CUDA_SUCCESS) { + av_log(NULL, AV_LOG_ERROR, "Error pushing cuda context\n"); + ret = AVERROR_UNKNOWN; + goto fail; + } + + err = cuModuleLoadData(&s->cu_module, vf_scale_cuda_ptx); + if (err != CUDA_SUCCESS) { + av_log(NULL, AV_LOG_ERROR, "Error loading module data\n"); + ret = AVERROR_UNKNOWN; + goto fail; + } + + cuModuleGetFunction(&s->cu_func_uchar, s->cu_module, "Subsample_Bilinear_uchar"); + cuModuleGetFunction(&s->cu_func_uchar2, s->cu_module, "Subsample_Bilinear_uchar2"); + cuModuleGetFunction(&s->cu_func_uchar4, s->cu_module, "Subsample_Bilinear_uchar4"); + cuModuleGetFunction(&s->cu_func_ushort, s->cu_module, "Subsample_Bilinear_ushort"); + cuModuleGetFunction(&s->cu_func_ushort2, s->cu_module, "Subsample_Bilinear_ushort2"); + cuModuleGetFunction(&s->cu_func_ushort4, s->cu_module, "Subsample_Bilinear_ushort4"); + + cuModuleGetTexRef(&s->cu_tex_uchar, s->cu_module, "uchar_tex"); + cuModuleGetTexRef(&s->cu_tex_uchar2, s->cu_module, "uchar2_tex"); + cuModuleGetTexRef(&s->cu_tex_uchar4, s->cu_module, "uchar4_tex"); + cuModuleGetTexRef(&s->cu_tex_ushort, s->cu_module, "ushort_tex"); + cuModuleGetTexRef(&s->cu_tex_ushort2, s->cu_module, "ushort2_tex"); + cuModuleGetTexRef(&s->cu_tex_ushort4, s->cu_module, "ushort4_tex"); + + cuTexRefSetFlags(s->cu_tex_uchar, CU_TRSF_READ_AS_INTEGER); + cuTexRefSetFlags(s->cu_tex_uchar2, CU_TRSF_READ_AS_INTEGER); + cuTexRefSetFlags(s->cu_tex_uchar4, CU_TRSF_READ_AS_INTEGER); + cuTexRefSetFlags(s->cu_tex_ushort, CU_TRSF_READ_AS_INTEGER); + cuTexRefSetFlags(s->cu_tex_ushort2, CU_TRSF_READ_AS_INTEGER); + cuTexRefSetFlags(s->cu_tex_ushort4, CU_TRSF_READ_AS_INTEGER); + + cuTexRefSetFilterMode(s->cu_tex_uchar, CU_TR_FILTER_MODE_LINEAR); + cuTexRefSetFilterMode(s->cu_tex_uchar2, CU_TR_FILTER_MODE_LINEAR); + cuTexRefSetFilterMode(s->cu_tex_uchar4, CU_TR_FILTER_MODE_LINEAR); + cuTexRefSetFilterMode(s->cu_tex_ushort, CU_TR_FILTER_MODE_LINEAR); + cuTexRefSetFilterMode(s->cu_tex_ushort2, CU_TR_FILTER_MODE_LINEAR); + cuTexRefSetFilterMode(s->cu_tex_ushort4, CU_TR_FILTER_MODE_LINEAR); + + cuCtxPopCurrent(&dummy); + + if ((ret = ff_scale_eval_dimensions(s, + s->w_expr, s->h_expr, + inlink, outlink, + &w, &h)) < 0) + goto fail; + + if (((int64_t)h * inlink->w) > INT_MAX || + ((int64_t)w * inlink->h) > INT_MAX) + av_log(ctx, AV_LOG_ERROR, "Rescaled value for width or height is too big.\n"); + + outlink->w = w; + outlink->h = h; + + ret = init_processing_chain(ctx, inlink->w, inlink->h, w, h); + if (ret < 0) + return ret; + + av_log(ctx, AV_LOG_VERBOSE, "w:%d h:%d -> w:%d h:%d\n", + inlink->w, inlink->h, outlink->w, outlink->h); + + if (inlink->sample_aspect_ratio.num) + outlink->sample_aspect_ratio = av_mul_q((AVRational){outlink->h*inlink->w, + outlink->w*inlink->h}, + inlink->sample_aspect_ratio); + else + outlink->sample_aspect_ratio = inlink->sample_aspect_ratio; + + return 0; + +fail: + return ret; +} + +static int call_resize_kernel(CUDAScaleContext *s, CUfunction func, CUtexref tex, int channels, + uint8_t *src_dptr, int src_width, int src_height, int src_pitch, + uint8_t *dst_dptr, int dst_width, int dst_height, int dst_pitch, + int pixel_size) +{ + CUdeviceptr src_devptr = (CUdeviceptr)src_dptr; + CUdeviceptr dst_devptr = (CUdeviceptr)dst_dptr; + void *args_uchar[] = { &dst_devptr, &dst_width, &dst_height, &dst_pitch, &src_width, &src_height }; + CUDA_ARRAY_DESCRIPTOR desc; + + desc.Width = src_width; + desc.Height = src_height; + desc.NumChannels = channels; + if (pixel_size == 1) { + desc.Format = CU_AD_FORMAT_UNSIGNED_INT8; + } else { + desc.Format = CU_AD_FORMAT_UNSIGNED_INT16; + } + + cuTexRefSetAddress2D_v3(tex, &desc, src_devptr, src_pitch * pixel_size); + cuLaunchKernel(func, DIV_UP(dst_width, BLOCKX), DIV_UP(dst_height, BLOCKY), 1, BLOCKX, BLOCKY, 1, 0, 0, args_uchar, NULL); + + return 0; +} + +static int scalecuda_resize(AVFilterContext *ctx, + AVFrame *out, AVFrame *in) +{ + AVHWFramesContext *in_frames_ctx = (AVHWFramesContext*)in->hw_frames_ctx->data; + CUDAScaleContext *s = ctx->priv; + + switch (in_frames_ctx->sw_format) { + case AV_PIX_FMT_YUV420P: + call_resize_kernel(s, s->cu_func_uchar, s->cu_tex_uchar, 1, + in->data[0], in->width, in->height, in->linesize[0], + out->data[0], out->width, out->height, out->linesize[0], + 1); + call_resize_kernel(s, s->cu_func_uchar, s->cu_tex_uchar, 1, + in->data[0]+in->linesize[0]*in->height, in->width/2, in->height/2, in->linesize[0]/2, + out->data[0]+out->linesize[0]*out->height, out->width/2, out->height/2, out->linesize[0]/2, + 1); + call_resize_kernel(s, s->cu_func_uchar, s->cu_tex_uchar, 1, + in->data[0]+ ALIGN_UP((in->linesize[0]*in->height*5)/4, s->tex_alignment), in->width/2, in->height/2, in->linesize[0]/2, + out->data[0]+(out->linesize[0]*out->height*5)/4, out->width/2, out->height/2, out->linesize[0]/2, + 1); + break; + case AV_PIX_FMT_YUV444P: + call_resize_kernel(s, s->cu_func_uchar, s->cu_tex_uchar, 1, + in->data[0], in->width, in->height, in->linesize[0], + out->data[0], out->width, out->height, out->linesize[0], + 1); + call_resize_kernel(s, s->cu_func_uchar, s->cu_tex_uchar, 1, + in->data[0]+in->linesize[0]*in->height, in->width, in->height, in->linesize[0], + out->data[0]+out->linesize[0]*out->height, out->width, out->height, out->linesize[0], + 1); + call_resize_kernel(s, s->cu_func_uchar, s->cu_tex_uchar, 1, + in->data[0]+in->linesize[0]*in->height*2, in->width, in->height, in->linesize[0], + out->data[0]+out->linesize[0]*out->height*2, out->width, out->height, out->linesize[0], + 1); + break; + case AV_PIX_FMT_NV12: + call_resize_kernel(s, s->cu_func_uchar, s->cu_tex_uchar, 1, + in->data[0], in->width, in->height, in->linesize[0], + out->data[0], out->width, out->height, out->linesize[0], + 1); + call_resize_kernel(s, s->cu_func_uchar2, s->cu_tex_uchar2, 2, + in->data[1], in->width/2, in->height/2, in->linesize[1], + out->data[0] + out->linesize[0] * ((out->height + 31) & ~0x1f), out->width/2, out->height/2, out->linesize[1]/2, + 1); + break; + case AV_PIX_FMT_P010LE: + call_resize_kernel(s, s->cu_func_ushort, s->cu_tex_ushort, 1, + in->data[0], in->width, in->height, in->linesize[0]/2, + out->data[0], out->width, out->height, out->linesize[0]/2, + 2); + call_resize_kernel(s, s->cu_func_ushort2, s->cu_tex_ushort2, 2, + in->data[1], in->width / 2, in->height / 2, in->linesize[1]/2, + out->data[0] + out->linesize[0] * ((out->height + 31) & ~0x1f), out->width / 2, out->height / 2, out->linesize[1] / 4, + 2); + break; + case AV_PIX_FMT_P016LE: + call_resize_kernel(s, s->cu_func_ushort, s->cu_tex_ushort, 1, + in->data[0], in->width, in->height, in->linesize[0] / 2, + out->data[0], out->width, out->height, out->linesize[0] / 2, + 2); + call_resize_kernel(s, s->cu_func_ushort2, s->cu_tex_ushort2, 2, + in->data[1], in->width / 2, in->height / 2, in->linesize[1] / 2, + out->data[0] + out->linesize[0] * ((out->height + 31) & ~0x1f), out->width / 2, out->height / 2, out->linesize[1] / 4, + 2); + break; + default: + return AVERROR_BUG; + } + + return 0; +} + +static int cudascale_scale(AVFilterContext *ctx, AVFrame *out, AVFrame *in) +{ + CUDAScaleContext *s = ctx->priv; + AVFrame *src = in; + int ret; + + ret = scalecuda_resize(ctx, s->frame, src); + if (ret < 0) + return ret; + + src = s->frame; + ret = av_hwframe_get_buffer(src->hw_frames_ctx, s->tmp_frame, 0); + if (ret < 0) + return ret; + + av_frame_move_ref(out, s->frame); + av_frame_move_ref(s->frame, s->tmp_frame); + + ret = av_frame_copy_props(out, in); + if (ret < 0) + return ret; + + return 0; +} + +static int cudascale_filter_frame(AVFilterLink *link, AVFrame *in) +{ + AVFilterContext *ctx = link->dst; + CUDAScaleContext *s = ctx->priv; + AVFilterLink *outlink = ctx->outputs[0]; + AVHWFramesContext *frames_ctx = (AVHWFramesContext*)s->frames_ctx->data; + AVCUDADeviceContext *device_hwctx = frames_ctx->device_ctx->hwctx; + + AVFrame *out = NULL; + CUresult err; + CUcontext dummy; + int ret = 0; + + out = av_frame_alloc(); + if (!out) { + ret = AVERROR(ENOMEM); + goto fail; + } + + err = cuCtxPushCurrent(device_hwctx->cuda_ctx); + if (err != CUDA_SUCCESS) { + ret = AVERROR_UNKNOWN; + goto fail; + } + + ret = cudascale_scale(ctx, out, in); + + cuCtxPopCurrent(&dummy); + if (ret < 0) + goto fail; + + av_reduce(&out->sample_aspect_ratio.num, &out->sample_aspect_ratio.den, + (int64_t)in->sample_aspect_ratio.num * outlink->h * link->w, + (int64_t)in->sample_aspect_ratio.den * outlink->w * link->h, + INT_MAX); + + av_frame_free(&in); + return ff_filter_frame(outlink, out); +fail: + av_frame_free(&in); + av_frame_free(&out); + return ret; +} + +#define OFFSET(x) offsetof(CUDAScaleContext, x) +#define FLAGS (AV_OPT_FLAG_FILTERING_PARAM|AV_OPT_FLAG_VIDEO_PARAM) +static const AVOption options[] = { + { "w", "Output video width", OFFSET(w_expr), AV_OPT_TYPE_STRING, { .str = "iw" }, .flags = FLAGS }, + { "h", "Output video height", OFFSET(h_expr), AV_OPT_TYPE_STRING, { .str = "ih" }, .flags = FLAGS }, + { NULL }, +}; + +static const AVClass cudascale_class = { + .class_name = "cudascale", + .item_name = av_default_item_name, + .option = options, + .version = LIBAVUTIL_VERSION_INT, +}; + +static const AVFilterPad cudascale_inputs[] = { + { + .name = "default", + .type = AVMEDIA_TYPE_VIDEO, + .filter_frame = cudascale_filter_frame, + }, + { NULL } +}; + +static const AVFilterPad cudascale_outputs[] = { + { + .name = "default", + .type = AVMEDIA_TYPE_VIDEO, + .config_props = cudascale_config_props, + }, + { NULL } +}; + +AVFilter ff_vf_scale_cuda = { + .name = "scale_cuda", + .description = NULL_IF_CONFIG_SMALL("GPU accelerated video resizer"), + + .init = cudascale_init, + .uninit = cudascale_uninit, + .query_formats = cudascale_query_formats, + + .priv_size = sizeof(CUDAScaleContext), + .priv_class = &cudascale_class, + + .inputs = cudascale_inputs, + .outputs = cudascale_outputs, + + .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE, +}; diff --git a/libavfilter/vf_scale_cuda.cu b/libavfilter/vf_scale_cuda.cu new file mode 100644 index 0000000000..5f5ec81989 --- /dev/null +++ b/libavfilter/vf_scale_cuda.cu @@ -0,0 +1,212 @@ +/* + * Copyright (c) 2017, NVIDIA CORPORATION. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a + * copy of this software and associated documentation files (the "Software"), + * to deal in the Software without restriction, including without limitation + * the rights to use, copy, modify, merge, publish, distribute, sublicense, + * and/or sell copies of the Software, and to permit persons to whom the + * Software is furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL + * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER + * DEALINGS IN THE SOFTWARE. + */ + +extern "C" { + +texture uchar_tex; +texture uchar2_tex; +texture uchar4_tex; +texture ushort_tex; +texture ushort2_tex; +texture ushort4_tex; + +__global__ void Subsample_Bilinear_uchar(unsigned char *dst, + int dst_width, int dst_height, int dst_pitch, + int src_width, int src_height) +{ + int xo = blockIdx.x * blockDim.x + threadIdx.x; + int yo = blockIdx.y * blockDim.y + threadIdx.y; + + if (yo < dst_height && xo < dst_width) + { + float hscale = (float)src_width / (float)dst_width; + float vscale = (float)src_height / (float)dst_height; + float xi = (xo + 0.5f) * hscale; + float yi = (yo + 0.5f) * vscale; + // 3-tap filter weights are {wh,1.0,wh} and {wv,1.0,wv} + float wh = min(max(0.5f * (hscale - 1.0f), 0.0f), 1.0f); + float wv = min(max(0.5f * (vscale - 1.0f), 0.0f), 1.0f); + // Convert weights to two bilinear weights -> {wh,1.0,wh} -> {wh,0.5,0} + {0,0.5,wh} + float dx = wh / (0.5f + wh); + float dy = wv / (0.5f + wv); + int y0 = tex2D(uchar_tex, xi-dx, yi-dy); + int y1 = tex2D(uchar_tex, xi+dx, yi-dy); + int y2 = tex2D(uchar_tex, xi-dx, yi+dy); + int y3 = tex2D(uchar_tex, xi+dx, yi+dy); + dst[yo*dst_pitch+xo] = (unsigned char)((y0+y1+y2+y3+2) >> 2); + } +} + +__global__ void Subsample_Bilinear_uchar2(uchar2 *dst, + int dst_width, int dst_height, int dst_pitch2, + int src_width, int src_height) +{ + int xo = blockIdx.x * blockDim.x + threadIdx.x; + int yo = blockIdx.y * blockDim.y + threadIdx.y; + + if (yo < dst_height && xo < dst_width) + { + float hscale = (float)src_width / (float)dst_width; + float vscale = (float)src_height / (float)dst_height; + float xi = (xo + 0.5f) * hscale; + float yi = (yo + 0.5f) * vscale; + // 3-tap filter weights are {wh,1.0,wh} and {wv,1.0,wv} + float wh = min(max(0.5f * (hscale - 1.0f), 0.0f), 1.0f); + float wv = min(max(0.5f * (vscale - 1.0f), 0.0f), 1.0f); + // Convert weights to two bilinear weights -> {wh,1.0,wh} -> {wh,0.5,0} + {0,0.5,wh} + float dx = wh / (0.5f + wh); + float dy = wv / (0.5f + wv); + uchar2 c0 = tex2D(uchar2_tex, xi-dx, yi-dy); + uchar2 c1 = tex2D(uchar2_tex, xi+dx, yi-dy); + uchar2 c2 = tex2D(uchar2_tex, xi-dx, yi+dy); + uchar2 c3 = tex2D(uchar2_tex, xi+dx, yi+dy); + int2 uv; + uv.x = ((int)c0.x+(int)c1.x+(int)c2.x+(int)c3.x+2) >> 2; + uv.y = ((int)c0.y+(int)c1.y+(int)c2.y+(int)c3.y+2) >> 2; + dst[yo*dst_pitch2+xo] = make_uchar2((unsigned char)uv.x, (unsigned char)uv.y); + } +} + +__global__ void Subsample_Bilinear_uchar4(uchar4 *dst, + int dst_width, int dst_height, int dst_pitch, + int src_width, int src_height) +{ + int xo = blockIdx.x * blockDim.x + threadIdx.x; + int yo = blockIdx.y * blockDim.y + threadIdx.y; + + if (yo < dst_height && xo < dst_width) + { + float hscale = (float)src_width / (float)dst_width; + float vscale = (float)src_height / (float)dst_height; + float xi = (xo + 0.5f) * hscale; + float yi = (yo + 0.5f) * vscale; + // 3-tap filter weights are {wh,1.0,wh} and {wv,1.0,wv} + float wh = min(max(0.5f * (hscale - 1.0f), 0.0f), 1.0f); + float wv = min(max(0.5f * (vscale - 1.0f), 0.0f), 1.0f); + // Convert weights to two bilinear weights -> {wh,1.0,wh} -> {wh,0.5,0} + {0,0.5,wh} + float dx = wh / (0.5f + wh); + float dy = wv / (0.5f + wv); + uchar4 c0 = tex2D(uchar4_tex, xi-dx, yi-dy); + uchar4 c1 = tex2D(uchar4_tex, xi+dx, yi-dy); + uchar4 c2 = tex2D(uchar4_tex, xi-dx, yi+dy); + uchar4 c3 = tex2D(uchar4_tex, xi+dx, yi+dy); + int4 res; + res.x = ((int)c0.x+(int)c1.x+(int)c2.x+(int)c3.x+2) >> 2; + res.y = ((int)c0.y+(int)c1.y+(int)c2.y+(int)c3.y+2) >> 2; + res.z = ((int)c0.z+(int)c1.z+(int)c2.z+(int)c3.z+2) >> 2; + res.w = ((int)c0.w+(int)c1.w+(int)c2.w+(int)c3.w+2) >> 2; + dst[yo*dst_pitch+xo] = make_uchar4( + (unsigned char)res.x, (unsigned char)res.y, (unsigned char)res.z, (unsigned char)res.w); + } +} + +__global__ void Subsample_Bilinear_ushort(unsigned short *dst, + int dst_width, int dst_height, int dst_pitch, + int src_width, int src_height) +{ + int xo = blockIdx.x * blockDim.x + threadIdx.x; + int yo = blockIdx.y * blockDim.y + threadIdx.y; + + if (yo < dst_height && xo < dst_width) + { + float hscale = (float)src_width / (float)dst_width; + float vscale = (float)src_height / (float)dst_height; + float xi = (xo + 0.5f) * hscale; + float yi = (yo + 0.5f) * vscale; + // 3-tap filter weights are {wh,1.0,wh} and {wv,1.0,wv} + float wh = min(max(0.5f * (hscale - 1.0f), 0.0f), 1.0f); + float wv = min(max(0.5f * (vscale - 1.0f), 0.0f), 1.0f); + // Convert weights to two bilinear weights -> {wh,1.0,wh} -> {wh,0.5,0} + {0,0.5,wh} + float dx = wh / (0.5f + wh); + float dy = wv / (0.5f + wv); + int y0 = tex2D(ushort_tex, xi-dx, yi-dy); + int y1 = tex2D(ushort_tex, xi+dx, yi-dy); + int y2 = tex2D(ushort_tex, xi-dx, yi+dy); + int y3 = tex2D(ushort_tex, xi+dx, yi+dy); + dst[yo*dst_pitch+xo] = (unsigned short)((y0+y1+y2+y3+2) >> 2); + } +} + +__global__ void Subsample_Bilinear_ushort2(ushort2 *dst, + int dst_width, int dst_height, int dst_pitch2, + int src_width, int src_height) +{ + int xo = blockIdx.x * blockDim.x + threadIdx.x; + int yo = blockIdx.y * blockDim.y + threadIdx.y; + + if (yo < dst_height && xo < dst_width) + { + float hscale = (float)src_width / (float)dst_width; + float vscale = (float)src_height / (float)dst_height; + float xi = (xo + 0.5f) * hscale; + float yi = (yo + 0.5f) * vscale; + // 3-tap filter weights are {wh,1.0,wh} and {wv,1.0,wv} + float wh = min(max(0.5f * (hscale - 1.0f), 0.0f), 1.0f); + float wv = min(max(0.5f * (vscale - 1.0f), 0.0f), 1.0f); + // Convert weights to two bilinear weights -> {wh,1.0,wh} -> {wh,0.5,0} + {0,0.5,wh} + float dx = wh / (0.5f + wh); + float dy = wv / (0.5f + wv); + ushort2 c0 = tex2D(ushort2_tex, xi-dx, yi-dy); + ushort2 c1 = tex2D(ushort2_tex, xi+dx, yi-dy); + ushort2 c2 = tex2D(ushort2_tex, xi-dx, yi+dy); + ushort2 c3 = tex2D(ushort2_tex, xi+dx, yi+dy); + int2 uv; + uv.x = ((int)c0.x+(int)c1.x+(int)c2.x+(int)c3.x+2) >> 2; + uv.y = ((int)c0.y+(int)c1.y+(int)c2.y+(int)c3.y+2) >> 2; + dst[yo*dst_pitch2+xo] = make_ushort2((unsigned short)uv.x, (unsigned short)uv.y); + } +} + +__global__ void Subsample_Bilinear_ushort4(ushort4 *dst, + int dst_width, int dst_height, int dst_pitch, + int src_width, int src_height) +{ + int xo = blockIdx.x * blockDim.x + threadIdx.x; + int yo = blockIdx.y * blockDim.y + threadIdx.y; + + if (yo < dst_height && xo < dst_width) + { + float hscale = (float)src_width / (float)dst_width; + float vscale = (float)src_height / (float)dst_height; + float xi = (xo + 0.5f) * hscale; + float yi = (yo + 0.5f) * vscale; + // 3-tap filter weights are {wh,1.0,wh} and {wv,1.0,wv} + float wh = min(max(0.5f * (hscale - 1.0f), 0.0f), 1.0f); + float wv = min(max(0.5f * (vscale - 1.0f), 0.0f), 1.0f); + // Convert weights to two bilinear weights -> {wh,1.0,wh} -> {wh,0.5,0} + {0,0.5,wh} + float dx = wh / (0.5f + wh); + float dy = wv / (0.5f + wv); + ushort4 c0 = tex2D(ushort4_tex, xi-dx, yi-dy); + ushort4 c1 = tex2D(ushort4_tex, xi+dx, yi-dy); + ushort4 c2 = tex2D(ushort4_tex, xi-dx, yi+dy); + ushort4 c3 = tex2D(ushort4_tex, xi+dx, yi+dy); + int4 res; + res.x = ((int)c0.x+(int)c1.x+(int)c2.x+(int)c3.x+2) >> 2; + res.y = ((int)c0.y+(int)c1.y+(int)c2.y+(int)c3.y+2) >> 2; + res.z = ((int)c0.z+(int)c1.z+(int)c2.z+(int)c3.z+2) >> 2; + res.w = ((int)c0.w+(int)c1.w+(int)c2.w+(int)c3.w+2) >> 2; + dst[yo*dst_pitch+xo] = make_ushort4( + (unsigned short)res.x, (unsigned short)res.y, (unsigned short)res.z, (unsigned short)res.w); + } +} + +}