virtualenv/FFmpeg
1
0
Fork 0
mirror of https://github.com/FFmpeg/FFmpeg.git synced 2025-08-10 06:10:52 +02:00
Code Issues Releases Activity

avfilter/vf_yadif_cuda: Switch to using ffnvcodec

Browse Source 
This change switches the vf_thumbnail_cuda filter from using the
full cuda sdk to using the ffnvcodec headers and loader.

Signed-off-by: Philip Langdale <philipl@overt.org>
Signed-off-by: Timo Rothenpieler <timo@rothenpieler.org>
This commit is contained in:
Philip Langdale
2019-02-20 19:57:50 -08:00
committed by Timo Rothenpieler
parent 5f47bfd50a
commit 7debf4277c
2 changed files with 32 additions and 28 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
configure vendored
View File

@@ -3535,7 +3535,7 @@ zscale_filter_deps="libzimg const_nan"
scale_vaapi_filter_deps="vaapi" scale_vaapi_filter_deps="vaapi"
vpp_qsv_filter_deps="libmfx" vpp_qsv_filter_deps="libmfx"
vpp_qsv_filter_select="qsvvpp" vpp_qsv_filter_select="qsvvpp"
yadif_cuda_filter_deps="cuda_sdk" yadif_cuda_filter_deps="ffnvcodec cuda_nvcc"
# examples # examples
avio_dir_cmd_deps="avformat avutil" avio_dir_cmd_deps="avformat avutil"

58
libavfilter/vf_yadif_cuda.c
View File

@@ -18,9 +18,8 @@
* Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
*/ */
#include <cuda.h>
#include "libavutil/avassert.h" #include "libavutil/avassert.h"
#include "libavutil/hwcontext_cuda.h" #include "libavutil/hwcontext_cuda_internal.h"
#include "libavutil/cuda_check.h" #include "libavutil/cuda_check.h"
#include "internal.h" #include "internal.h"
#include "yadif.h" #include "yadif.h"
@@ -49,7 +48,7 @@ typedef struct DeintCUDAContext {
#define BLOCKX 32 #define BLOCKX 32
#define BLOCKY 16 #define BLOCKY 16
#define CHECK_CU(x) FF_CUDA_CHECK(ctx, x) #define CHECK_CU(x) FF_CUDA_CHECK_DL(ctx, s->hwctx->internal->cuda_dl, x)
static CUresult call_kernel(AVFilterContext *ctx, CUfunction func, static CUresult call_kernel(AVFilterContext *ctx, CUfunction func,
CUdeviceptr prev, CUdeviceptr cur, CUdeviceptr next, CUdeviceptr prev, CUdeviceptr cur, CUdeviceptr next,
@@ -64,6 +63,7 @@ static CUresult call_kernel(AVFilterContext *ctx, CUfunction func,
int parity, int tff) int parity, int tff)
{ {
DeintCUDAContext *s = ctx->priv; DeintCUDAContext *s = ctx->priv;
CudaFunctions *cu = s->hwctx->internal->cuda_dl;
CUtexObject tex_prev = 0, tex_cur = 0, tex_next = 0; CUtexObject tex_prev = 0, tex_cur = 0, tex_next = 0;
int ret; int ret;
int skip_spatial_check = s->yadif.mode&2; int skip_spatial_check = s->yadif.mode&2;
@@ -88,32 +88,32 @@ static CUresult call_kernel(AVFilterContext *ctx, CUfunction func,
}; };
res_desc.res.pitch2D.devPtr = (CUdeviceptr)prev; res_desc.res.pitch2D.devPtr = (CUdeviceptr)prev;
ret = CHECK_CU(cuTexObjectCreate(&tex_prev, &res_desc, &tex_desc, NULL)); ret = CHECK_CU(cu->cuTexObjectCreate(&tex_prev, &res_desc, &tex_desc, NULL));
if (ret < 0) if (ret < 0)
goto exit; goto exit;
res_desc.res.pitch2D.devPtr = (CUdeviceptr)cur; res_desc.res.pitch2D.devPtr = (CUdeviceptr)cur;
ret = CHECK_CU(cuTexObjectCreate(&tex_cur, &res_desc, &tex_desc, NULL)); ret = CHECK_CU(cu->cuTexObjectCreate(&tex_cur, &res_desc, &tex_desc, NULL));
if (ret < 0) if (ret < 0)
goto exit; goto exit;
res_desc.res.pitch2D.devPtr = (CUdeviceptr)next; res_desc.res.pitch2D.devPtr = (CUdeviceptr)next;
ret = CHECK_CU(cuTexObjectCreate(&tex_next, &res_desc, &tex_desc, NULL)); ret = CHECK_CU(cu->cuTexObjectCreate(&tex_next, &res_desc, &tex_desc, NULL));
if (ret < 0) if (ret < 0)
goto exit; goto exit;
ret = CHECK_CU(cuLaunchKernel(func, ret = CHECK_CU(cu->cuLaunchKernel(func,
DIV_UP(dst_width, BLOCKX), DIV_UP(dst_height, BLOCKY), 1, DIV_UP(dst_width, BLOCKX), DIV_UP(dst_height, BLOCKY), 1,
BLOCKX, BLOCKY, 1, BLOCKX, BLOCKY, 1,
0, s->stream, args, NULL)); 0, s->stream, args, NULL));
exit: exit:
if (tex_prev) if (tex_prev)
CHECK_CU(cuTexObjectDestroy(tex_prev)); CHECK_CU(cu->cuTexObjectDestroy(tex_prev));
if (tex_cur) if (tex_cur)
CHECK_CU(cuTexObjectDestroy(tex_cur)); CHECK_CU(cu->cuTexObjectDestroy(tex_cur));
if (tex_next) if (tex_next)
CHECK_CU(cuTexObjectDestroy(tex_next)); CHECK_CU(cu->cuTexObjectDestroy(tex_next));
return ret; return ret;
} }
@@ -123,10 +123,11 @@ static void filter(AVFilterContext *ctx, AVFrame *dst,
{ {
DeintCUDAContext *s = ctx->priv; DeintCUDAContext *s = ctx->priv;
YADIFContext *y = &s->yadif; YADIFContext *y = &s->yadif;
CudaFunctions *cu = s->hwctx->internal->cuda_dl;
CUcontext dummy; CUcontext dummy;
int i, ret; int i, ret;
ret = CHECK_CU(cuCtxPushCurrent(s->cu_ctx)); ret = CHECK_CU(cu->cuCtxPushCurrent(s->cu_ctx));
if (ret < 0) if (ret < 0)
return; return;
@@ -179,10 +180,10 @@ static void filter(AVFilterContext *ctx, AVFrame *dst,
parity, tff); parity, tff);
} }
CHECK_CU(cuStreamSynchronize(s->stream)); CHECK_CU(cu->cuStreamSynchronize(s->stream));
exit: exit:
CHECK_CU(cuCtxPopCurrent(&dummy)); CHECK_CU(cu->cuCtxPopCurrent(&dummy));
return; return;
} }
@@ -192,10 +193,11 @@ static av_cold void deint_cuda_uninit(AVFilterContext *ctx)
DeintCUDAContext *s = ctx->priv; DeintCUDAContext *s = ctx->priv;
YADIFContext *y = &s->yadif; YADIFContext *y = &s->yadif;
if (s->cu_module) { if (s->hwctx && s->cu_module) {
CHECK_CU(cuCtxPushCurrent(s->cu_ctx)); CudaFunctions *cu = s->hwctx->internal->cuda_dl;
CHECK_CU(cuModuleUnload(s->cu_module)); CHECK_CU(cu->cuCtxPushCurrent(s->cu_ctx));
CHECK_CU(cuCtxPopCurrent(&dummy)); CHECK_CU(cu->cuModuleUnload(s->cu_module));
CHECK_CU(cu->cuCtxPopCurrent(&dummy));
} }
av_frame_free(&y->prev); av_frame_free(&y->prev);
@@ -253,6 +255,7 @@ static int config_output(AVFilterLink *link)
AVFilterContext *ctx = link->src; AVFilterContext *ctx = link->src;
DeintCUDAContext *s = ctx->priv; DeintCUDAContext *s = ctx->priv;
YADIFContext *y = &s->yadif; YADIFContext *y = &s->yadif;
CudaFunctions *cu;
int ret = 0; int ret = 0;
CUcontext dummy; CUcontext dummy;
@@ -266,6 +269,7 @@ static int config_output(AVFilterLink *link)
s->hwctx = ((AVHWDeviceContext*)s->device_ref->data)->hwctx; s->hwctx = ((AVHWDeviceContext*)s->device_ref->data)->hwctx;
s->cu_ctx = s->hwctx->cuda_ctx; s->cu_ctx = s->hwctx->cuda_ctx;
s->stream = s->hwctx->stream; s->stream = s->hwctx->stream;
cu = s->hwctx->internal->cuda_dl;
link->hw_frames_ctx = av_hwframe_ctx_alloc(s->device_ref); link->hw_frames_ctx = av_hwframe_ctx_alloc(s->device_ref);
if (!link->hw_frames_ctx) { if (!link->hw_frames_ctx) {
@@ -313,32 +317,32 @@ static int config_output(AVFilterLink *link)
y->csp = av_pix_fmt_desc_get(output_frames->sw_format); y->csp = av_pix_fmt_desc_get(output_frames->sw_format);
y->filter = filter; y->filter = filter;
ret = CHECK_CU(cuCtxPushCurrent(s->cu_ctx)); ret = CHECK_CU(cu->cuCtxPushCurrent(s->cu_ctx));
if (ret < 0) if (ret < 0)
goto exit; goto exit;
ret = CHECK_CU(cuModuleLoadData(&s->cu_module, vf_yadif_cuda_ptx)); ret = CHECK_CU(cu->cuModuleLoadData(&s->cu_module, vf_yadif_cuda_ptx));
if (ret < 0) if (ret < 0)
goto exit; goto exit;
ret = CHECK_CU(cuModuleGetFunction(&s->cu_func_uchar, s->cu_module, "yadif_uchar")); ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar, s->cu_module, "yadif_uchar"));
if (ret < 0) if (ret < 0)
goto exit; goto exit;
ret = CHECK_CU(cuModuleGetFunction(&s->cu_func_uchar2, s->cu_module, "yadif_uchar2")); ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar2, s->cu_module, "yadif_uchar2"));
if (ret < 0) if (ret < 0)
goto exit; goto exit;
ret = CHECK_CU(cuModuleGetFunction(&s->cu_func_ushort, s->cu_module, "yadif_ushort")); ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_ushort, s->cu_module, "yadif_ushort"));
if (ret < 0) if (ret < 0)
goto exit; goto exit;
ret = CHECK_CU(cuModuleGetFunction(&s->cu_func_ushort2, s->cu_module, "yadif_ushort2")); ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_ushort2, s->cu_module, "yadif_ushort2"));
if (ret < 0) if (ret < 0)
goto exit; goto exit;
exit: exit:
CHECK_CU(cuCtxPopCurrent(&dummy)); CHECK_CU(cu->cuCtxPopCurrent(&dummy));
return ret; return ret;
} }