X-Git-Url: https://git.sesse.net/?a=blobdiff_plain;f=libavfilter%2Fvf_yadif_cuda.c;h=4e41c8b5546ddd69b8f1cd1ff6c7b3670cd1014e;hb=a04ad248a05e7b613abe09b3bb067f555108d794;hp=85e1aac5ebdebf3218510fc0e60a814cfeec031f;hpb=bec3b2041dcc20ab4b06a6b31d09130e1a7166c3;p=ffmpeg diff --git a/libavfilter/vf_yadif_cuda.c b/libavfilter/vf_yadif_cuda.c index 85e1aac5ebd..4e41c8b5546 100644 --- a/libavfilter/vf_yadif_cuda.c +++ b/libavfilter/vf_yadif_cuda.c @@ -18,9 +18,8 @@ * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA */ -#include #include "libavutil/avassert.h" -#include "libavutil/hwcontext_cuda.h" +#include "libavutil/hwcontext_cuda_internal.h" #include "libavutil/cuda_check.h" #include "internal.h" #include "yadif.h" @@ -49,7 +48,7 @@ typedef struct DeintCUDAContext { #define BLOCKX 32 #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, CUdeviceptr prev, CUdeviceptr cur, CUdeviceptr next, @@ -64,6 +63,7 @@ static CUresult call_kernel(AVFilterContext *ctx, CUfunction func, int parity, int tff) { DeintCUDAContext *s = ctx->priv; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; CUtexObject tex_prev = 0, tex_cur = 0, tex_next = 0; int ret; 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; - 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) goto exit; 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) goto exit; 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) goto exit; - ret = CHECK_CU(cuLaunchKernel(func, - DIV_UP(dst_width, BLOCKX), DIV_UP(dst_height, BLOCKY), 1, - BLOCKX, BLOCKY, 1, - 0, s->stream, args, NULL)); + ret = CHECK_CU(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)); + CHECK_CU(cu->cuTexObjectDestroy(tex_prev)); if (tex_cur) - CHECK_CU(cuTexObjectDestroy(tex_cur)); + CHECK_CU(cu->cuTexObjectDestroy(tex_cur)); if (tex_next) - CHECK_CU(cuTexObjectDestroy(tex_next)); + CHECK_CU(cu->cuTexObjectDestroy(tex_next)); return ret; } @@ -123,10 +123,11 @@ static void filter(AVFilterContext *ctx, AVFrame *dst, { DeintCUDAContext *s = ctx->priv; YADIFContext *y = &s->yadif; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; CUcontext dummy; int i, ret; - ret = CHECK_CU(cuCtxPushCurrent(s->cu_ctx)); + ret = CHECK_CU(cu->cuCtxPushCurrent(s->cu_ctx)); if (ret < 0) return; @@ -179,10 +180,8 @@ static void filter(AVFilterContext *ctx, AVFrame *dst, parity, tff); } - CHECK_CU(cuStreamSynchronize(s->stream)); - exit: - CHECK_CU(cuCtxPopCurrent(&dummy)); + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); return; } @@ -192,10 +191,11 @@ static av_cold void deint_cuda_uninit(AVFilterContext *ctx) 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)); + if (s->hwctx && s->cu_module) { + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + CHECK_CU(cu->cuCtxPushCurrent(s->cu_ctx)); + CHECK_CU(cu->cuModuleUnload(s->cu_module)); + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); } av_frame_free(&y->prev); @@ -216,10 +216,10 @@ static int deint_cuda_query_formats(AVFilterContext *ctx) int ret; if ((ret = ff_formats_ref(ff_make_format_list(pix_fmts), - &ctx->inputs[0]->out_formats)) < 0) + &ctx->inputs[0]->outcfg.formats)) < 0) return ret; if ((ret = ff_formats_ref(ff_make_format_list(pix_fmts), - &ctx->outputs[0]->in_formats)) < 0) + &ctx->outputs[0]->incfg.formats)) < 0) return ret; return 0; @@ -253,6 +253,7 @@ static int config_output(AVFilterLink *link) AVFilterContext *ctx = link->src; DeintCUDAContext *s = ctx->priv; YADIFContext *y = &s->yadif; + CudaFunctions *cu; int ret = 0; CUcontext dummy; @@ -266,6 +267,7 @@ static int config_output(AVFilterLink *link) s->hwctx = ((AVHWDeviceContext*)s->device_ref->data)->hwctx; s->cu_ctx = s->hwctx->cuda_ctx; s->stream = s->hwctx->stream; + cu = s->hwctx->internal->cuda_dl; link->hw_frames_ctx = av_hwframe_ctx_alloc(s->device_ref); if (!link->hw_frames_ctx) { @@ -313,32 +315,32 @@ static int config_output(AVFilterLink *link) y->csp = av_pix_fmt_desc_get(output_frames->sw_format); y->filter = filter; - ret = CHECK_CU(cuCtxPushCurrent(s->cu_ctx)); + ret = CHECK_CU(cu->cuCtxPushCurrent(s->cu_ctx)); if (ret < 0) 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) 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) 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) 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) 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) goto exit; exit: - CHECK_CU(cuCtxPopCurrent(&dummy)); + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); return ret; } @@ -371,7 +373,7 @@ static const AVFilterPad deint_cuda_outputs[] = { { NULL } }; -AVFilter ff_vf_yadif_cuda = { +const AVFilter ff_vf_yadif_cuda = { .name = "yadif_cuda", .description = NULL_IF_CONFIG_SMALL("Deinterlace CUDA frames"), .priv_size = sizeof(DeintCUDAContext),