* 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"
#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)
+#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,
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;
- CUresult err;
+ int ret;
int skip_spatial_check = s->yadif.mode&2;
void *args[] = { &dst, &tex_prev, &tex_cur, &tex_next,
};
res_desc.res.pitch2D.devPtr = (CUdeviceptr)prev;
- err = CHECK_CU(cuTexObjectCreate(&tex_prev, &res_desc, &tex_desc, NULL));
- if (err != CUDA_SUCCESS) {
+ ret = CHECK_CU(cu->cuTexObjectCreate(&tex_prev, &res_desc, &tex_desc, NULL));
+ if (ret < 0)
goto exit;
- }
res_desc.res.pitch2D.devPtr = (CUdeviceptr)cur;
- err = CHECK_CU(cuTexObjectCreate(&tex_cur, &res_desc, &tex_desc, NULL));
- if (err != CUDA_SUCCESS) {
+ ret = CHECK_CU(cu->cuTexObjectCreate(&tex_cur, &res_desc, &tex_desc, NULL));
+ if (ret < 0)
goto exit;
- }
res_desc.res.pitch2D.devPtr = (CUdeviceptr)next;
- err = CHECK_CU(cuTexObjectCreate(&tex_next, &res_desc, &tex_desc, NULL));
- if (err != CUDA_SUCCESS) {
+ ret = CHECK_CU(cu->cuTexObjectCreate(&tex_next, &res_desc, &tex_desc, NULL));
+ if (ret < 0)
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));
+ 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 err;
+ return ret;
}
static void filter(AVFilterContext *ctx, AVFrame *dst,
{
DeintCUDAContext *s = ctx->priv;
YADIFContext *y = &s->yadif;
+ CudaFunctions *cu = s->hwctx->internal->cuda_dl;
CUcontext dummy;
- CUresult err;
- int i;
+ int i, ret;
- err = CHECK_CU(cuCtxPushCurrent(s->cu_ctx));
- if (err != CUDA_SUCCESS) {
- goto exit;
- }
+ ret = CHECK_CU(cu->cuCtxPushCurrent(s->cu_ctx));
+ if (ret < 0)
+ return;
for (i = 0; i < y->csp->nb_components; i++) {
CUfunction func;
parity, tff);
}
- err = CHECK_CU(cuStreamSynchronize(s->stream));
- if (err != CUDA_SUCCESS) {
- goto exit;
- }
-
exit:
- CHECK_CU(cuCtxPopCurrent(&dummy));
+ CHECK_CU(cu->cuCtxPopCurrent(&dummy));
return;
}
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);
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;
AVFilterContext *ctx = link->src;
DeintCUDAContext *s = ctx->priv;
YADIFContext *y = &s->yadif;
+ CudaFunctions *cu;
int ret = 0;
CUcontext dummy;
- CUresult err;
av_assert0(s->input_frames);
s->device_ref = av_buffer_ref(s->input_frames->device_ref);
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) {
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;
+ ret = CHECK_CU(cu->cuCtxPushCurrent(s->cu_ctx));
+ if (ret < 0)
goto exit;
- }
- err = CHECK_CU(cuModuleLoadData(&s->cu_module, vf_yadif_cuda_ptx));
- if (err != CUDA_SUCCESS) {
- ret = AVERROR_INVALIDDATA;
+ ret = CHECK_CU(cu->cuModuleLoadData(&s->cu_module, vf_yadif_cuda_ptx));
+ if (ret < 0)
goto exit;
- }
- err = CHECK_CU(cuModuleGetFunction(&s->cu_func_uchar, s->cu_module, "yadif_uchar"));
- if (err != CUDA_SUCCESS) {
- ret = AVERROR_INVALIDDATA;
+ ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar, s->cu_module, "yadif_uchar"));
+ if (ret < 0)
goto exit;
- }
- err = CHECK_CU(cuModuleGetFunction(&s->cu_func_uchar2, s->cu_module, "yadif_uchar2"));
- if (err != CUDA_SUCCESS) {
- ret = AVERROR_INVALIDDATA;
+ ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar2, s->cu_module, "yadif_uchar2"));
+ if (ret < 0)
goto exit;
- }
- err= CHECK_CU(cuModuleGetFunction(&s->cu_func_ushort, s->cu_module, "yadif_ushort"));
- if (err != CUDA_SUCCESS) {
- ret = AVERROR_INVALIDDATA;
+ ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_ushort, s->cu_module, "yadif_ushort"));
+ if (ret < 0)
goto exit;
- }
- err = CHECK_CU(cuModuleGetFunction(&s->cu_func_ushort2, s->cu_module, "yadif_ushort2"));
- if (err != CUDA_SUCCESS) {
- ret = AVERROR_INVALIDDATA;
+ 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;
}
{ 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),