avfilter: Constify all AVFilters

[ffmpeg] / libavfilter / vf_yadif_cuda.c

diff --git a/libavfilter/vf_yadif_cuda.c b/libavfilter/vf_yadif_cuda.c
index 85e1aac5ebdebf3218510fc0e60a814cfeec031f..4e41c8b5546ddd69b8f1cd1ff6c7b3670cd1014e 100644 (file)
--- 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 <cuda.h>
 #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),