libavfilter/vf_yadif_cuda.c

   1 /*
   2  * Copyright (C) 2018 Philip Langdale <philipl@overt.org>
   3  *
   4  * This file is part of FFmpeg.
   5  *
   6  * FFmpeg is free software; you can redistribute it and/or
   7  * modify it under the terms of the GNU Lesser General Public
   8  * License as published by the Free Software Foundation; either
   9  * version 2.1 of the License, or (at your option) any later version.
  10  *
  11  * FFmpeg is distributed in the hope that it will be useful,
  12  * but WITHOUT ANY WARRANTY; without even the implied warranty of
  13  * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
  14  * Lesser General Public License for more details.
  15  *
  16  * You should have received a copy of the GNU Lesser General Public
  17  * License along with FFmpeg; if not, write to the Free Software
  18  * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
  19  */
  20
  21 #include "libavutil/avassert.h"
  22 #include "libavutil/hwcontext.h"
  23 #include "libavutil/hwcontext_cuda_internal.h"
  24 #include "libavutil/cuda_check.h"
  25
  26 #include "filters.h"
  27 #include "yadif.h"
  28
  29 #include "cuda/load_helper.h"
  30
  31 extern const unsigned char ff_vf_yadif_cuda_ptx_data[];
  32 extern const unsigned int ff_vf_yadif_cuda_ptx_len;
  33
  34 typedef struct DeintCUDAContext {
  35     YADIFContext yadif;
  36
  37     AVCUDADeviceContext *hwctx;
  38     AVBufferRef         *device_ref;
  39     AVBufferRef         *input_frames_ref;
  40     AVHWFramesContext   *input_frames;
  41
  42     CUmodule    cu_module;
  43     CUfunction  cu_func_uchar;
  44     CUfunction  cu_func_uchar2;
  45     CUfunction  cu_func_ushort;
  46     CUfunction  cu_func_ushort2;
  47 } DeintCUDAContext;
  48
  49 #define DIV_UP(a, b) ( ((a) + (b) - 1) / (b) )
  50 #define ALIGN_UP(a, b) (((a) + (b) - 1) & ~((b) - 1))
  51 #define BLOCKX 32
  52 #define BLOCKY 16
  53
  54 #define CHECK_CU(x) FF_CUDA_CHECK_DL(ctx, s->hwctx->internal->cuda_dl, x)
  55
  56 static CUresult call_kernel(AVFilterContext *ctx, CUfunction func,
  57                             CUdeviceptr prev, CUdeviceptr cur, CUdeviceptr next,
  58                             CUarray_format format, int channels,
  59                             int src_width,  // Width is pixels per channel
  60                             int src_height, // Height is pixels per channel
  61                             int src_pitch,  // Pitch is bytes
  62                             CUdeviceptr dst,
  63                             int dst_width,  // Width is pixels per channel
  64                             int dst_height, // Height is pixels per channel
  65                             int dst_pitch,  // Pitch is pixels per channel
  66                             int parity, int tff)
  67 {
  68     DeintCUDAContext *s = ctx->priv;
  69     CudaFunctions *cu = s->hwctx->internal->cuda_dl;
  70     CUtexObject tex_prev = 0, tex_cur = 0, tex_next = 0;
  71     int ret;
  72     int skip_spatial_check = s->yadif.mode&2;
  73
  74     void *args[] = { &dst, &tex_prev, &tex_cur, &tex_next,
  75                      &dst_width, &dst_height, &dst_pitch,
  76                      &src_width, &src_height, &parity, &tff,
  77                      &skip_spatial_check };
  78
  79     CUDA_TEXTURE_DESC tex_desc = {
  80         .filterMode = CU_TR_FILTER_MODE_POINT,
  81         .flags = CU_TRSF_READ_AS_INTEGER,
  82     };
  83
  84     CUDA_RESOURCE_DESC res_desc = {
  85         .resType = CU_RESOURCE_TYPE_PITCH2D,
  86         .res.pitch2D.format = format,
  87         .res.pitch2D.numChannels = channels,
  88         .res.pitch2D.width = src_width,
  89         .res.pitch2D.height = src_height,
  90         .res.pitch2D.pitchInBytes = src_pitch,
  91     };
  92
  93     res_desc.res.pitch2D.devPtr = (CUdeviceptr)prev;
  94     ret = CHECK_CU(cu->cuTexObjectCreate(&tex_prev, &res_desc, &tex_desc, NULL));
  95     if (ret < 0)
  96         goto exit;
  97
  98     res_desc.res.pitch2D.devPtr = (CUdeviceptr)cur;
  99     ret = CHECK_CU(cu->cuTexObjectCreate(&tex_cur, &res_desc, &tex_desc, NULL));
 100     if (ret < 0)
 101         goto exit;
 102
 103     res_desc.res.pitch2D.devPtr = (CUdeviceptr)next;
 104     ret = CHECK_CU(cu->cuTexObjectCreate(&tex_next, &res_desc, &tex_desc, NULL));
 105     if (ret < 0)
 106         goto exit;
 107
 108     ret = CHECK_CU(cu->cuLaunchKernel(func,
 109                                       DIV_UP(dst_width, BLOCKX), DIV_UP(dst_height, BLOCKY), 1,
 110                                       BLOCKX, BLOCKY, 1,
 111                                       0, s->hwctx->stream, args, NULL));
 112
 113 exit:
 114     if (tex_prev)
 115         CHECK_CU(cu->cuTexObjectDestroy(tex_prev));
 116     if (tex_cur)
 117         CHECK_CU(cu->cuTexObjectDestroy(tex_cur));
 118     if (tex_next)
 119         CHECK_CU(cu->cuTexObjectDestroy(tex_next));
 120
 121     return ret;
 122 }
 123
 124 static void filter(AVFilterContext *ctx, AVFrame *dst,
 125                    int parity, int tff)
 126 {
 127     DeintCUDAContext *s = ctx->priv;
 128     YADIFContext *y = &s->yadif;
 129     CudaFunctions *cu = s->hwctx->internal->cuda_dl;
 130     CUcontext dummy;
 131     int i, ret;
 132
 133     ret = CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx));
 134     if (ret < 0)
 135         return;
 136
 137     for (i = 0; i < y->csp->nb_components; i++) {
 138         CUfunction func;
 139         CUarray_format format;
 140         int pixel_size, channels;
 141         const AVComponentDescriptor *comp = &y->csp->comp[i];
 142
 143         if (comp->plane < i) {
 144             // We process planes as a whole, so don't reprocess
 145             // them for additional components
 146             continue;
 147         }
 148
 149         pixel_size = (comp->depth + comp->shift) / 8;
 150         channels = comp->step / pixel_size;
 151         if (pixel_size > 2 || channels > 2) {
 152             av_log(ctx, AV_LOG_ERROR, "Unsupported pixel format: %s\n", y->csp->name);
 153             goto exit;
 154         }
 155         switch (pixel_size) {
 156         case 1:
 157             func = channels == 1 ? s->cu_func_uchar : s->cu_func_uchar2;
 158             format = CU_AD_FORMAT_UNSIGNED_INT8;
 159             break;
 160         case 2:
 161             func = channels == 1 ? s->cu_func_ushort : s->cu_func_ushort2;
 162             format = CU_AD_FORMAT_UNSIGNED_INT16;
 163             break;
 164         default:
 165             av_log(ctx, AV_LOG_ERROR, "Unsupported pixel format: %s\n", y->csp->name);
 166             goto exit;
 167         }
 168         av_log(ctx, AV_LOG_TRACE,
 169                "Deinterlacing plane %d: pixel_size: %d channels: %d\n",
 170                comp->plane, pixel_size, channels);
 171         call_kernel(ctx, func,
 172                     (CUdeviceptr)y->prev->data[i],
 173                     (CUdeviceptr)y->cur->data[i],
 174                     (CUdeviceptr)y->next->data[i],
 175                     format, channels,
 176                     AV_CEIL_RSHIFT(y->cur->width, i ? y->csp->log2_chroma_w : 0),
 177                     AV_CEIL_RSHIFT(y->cur->height, i ? y->csp->log2_chroma_h : 0),
 178                     y->cur->linesize[i],
 179                     (CUdeviceptr)dst->data[i],
 180                     AV_CEIL_RSHIFT(dst->width, i ? y->csp->log2_chroma_w : 0),
 181                     AV_CEIL_RSHIFT(dst->height, i ? y->csp->log2_chroma_h : 0),
 182                     dst->linesize[i] / comp->step,
 183                     parity, tff);
 184     }
 185
 186 exit:
 187     CHECK_CU(cu->cuCtxPopCurrent(&dummy));
 188     return;
 189 }
 190
 191 static av_cold void deint_cuda_uninit(AVFilterContext *ctx)
 192 {
 193     CUcontext dummy;
 194     DeintCUDAContext *s = ctx->priv;
 195
 196     if (s->hwctx && s->cu_module) {
 197         CudaFunctions *cu = s->hwctx->internal->cuda_dl;
 198         CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx));
 199         CHECK_CU(cu->cuModuleUnload(s->cu_module));
 200         CHECK_CU(cu->cuCtxPopCurrent(&dummy));
 201     }
 202
 203     ff_yadif_uninit(ctx);
 204
 205     av_buffer_unref(&s->device_ref);
 206     s->hwctx = NULL;
 207     av_buffer_unref(&s->input_frames_ref);
 208     s->input_frames = NULL;
 209 }
 210
 211 static int config_input(AVFilterLink *inlink)
 212 {
 213     FilterLink        *l = ff_filter_link(inlink);
 214     AVFilterContext *ctx = inlink->dst;
 215     DeintCUDAContext *s  = ctx->priv;
 216
 217     if (!l->hw_frames_ctx) {
 218         av_log(ctx, AV_LOG_ERROR, "A hardware frames reference is "
 219                "required to associate the processing device.\n");
 220         return AVERROR(EINVAL);
 221     }
 222
 223     s->input_frames_ref = av_buffer_ref(l->hw_frames_ctx);
 224     if (!s->input_frames_ref) {
 225         av_log(ctx, AV_LOG_ERROR, "A input frames reference create "
 226                "failed.\n");
 227         return AVERROR(ENOMEM);
 228     }
 229     s->input_frames = (AVHWFramesContext*)s->input_frames_ref->data;
 230
 231     return 0;
 232 }
 233
 234 static int config_output(AVFilterLink *link)
 235 {
 236     FilterLink *l = ff_filter_link(link);
 237     AVHWFramesContext *output_frames;
 238     AVFilterContext *ctx = link->src;
 239     DeintCUDAContext *s = ctx->priv;
 240     YADIFContext *y = &s->yadif;
 241     CudaFunctions *cu;
 242     int ret = 0;
 243     CUcontext dummy;
 244
 245     av_assert0(s->input_frames);
 246     s->device_ref = av_buffer_ref(s->input_frames->device_ref);
 247     if (!s->device_ref) {
 248         av_log(ctx, AV_LOG_ERROR, "A device reference create "
 249                "failed.\n");
 250         return AVERROR(ENOMEM);
 251     }
 252     s->hwctx = ((AVHWDeviceContext*)s->device_ref->data)->hwctx;
 253     cu = s->hwctx->internal->cuda_dl;
 254
 255     l->hw_frames_ctx = av_hwframe_ctx_alloc(s->device_ref);
 256     if (!l->hw_frames_ctx) {
 257         av_log(ctx, AV_LOG_ERROR, "Failed to create HW frame context "
 258                "for output.\n");
 259         ret = AVERROR(ENOMEM);
 260         goto exit;
 261     }
 262
 263     output_frames = (AVHWFramesContext*)l->hw_frames_ctx->data;
 264
 265     output_frames->format    = AV_PIX_FMT_CUDA;
 266     output_frames->sw_format = s->input_frames->sw_format;
 267     output_frames->width     = ctx->inputs[0]->w;
 268     output_frames->height    = ctx->inputs[0]->h;
 269
 270     output_frames->initial_pool_size = 4;
 271
 272     ret = ff_filter_init_hw_frames(ctx, link, 10);
 273     if (ret < 0)
 274         goto exit;
 275
 276     ret = av_hwframe_ctx_init(l->hw_frames_ctx);
 277     if (ret < 0) {
 278         av_log(ctx, AV_LOG_ERROR, "Failed to initialise CUDA frame "
 279                "context for output: %d\n", ret);
 280         goto exit;
 281     }
 282
 283     ret = ff_yadif_config_output_common(link);
 284     if (ret < 0)
 285         goto exit;
 286
 287     y->csp = av_pix_fmt_desc_get(output_frames->sw_format);
 288     y->filter = filter;
 289
 290     ret = CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx));
 291     if (ret < 0)
 292         goto exit;
 293
 294     ret = ff_cuda_load_module(ctx, s->hwctx, &s->cu_module, ff_vf_yadif_cuda_ptx_data, ff_vf_yadif_cuda_ptx_len);
 295     if (ret < 0)
 296         goto exit;
 297
 298     ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar, s->cu_module, "yadif_uchar"));
 299     if (ret < 0)
 300         goto exit;
 301
 302     ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar2, s->cu_module, "yadif_uchar2"));
 303     if (ret < 0)
 304         goto exit;
 305
 306     ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_ushort, s->cu_module, "yadif_ushort"));
 307     if (ret < 0)
 308         goto exit;
 309
 310     ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_ushort2, s->cu_module, "yadif_ushort2"));
 311     if (ret < 0)
 312         goto exit;
 313
 314 exit:
 315     CHECK_CU(cu->cuCtxPopCurrent(&dummy));
 316
 317     return ret;
 318 }
 319
 320 static const AVClass yadif_cuda_class = {
 321     .class_name = "yadif_cuda",
 322     .item_name  = av_default_item_name,
 323     .option     = ff_yadif_options,
 324     .version    = LIBAVUTIL_VERSION_INT,
 325     .category   = AV_CLASS_CATEGORY_FILTER,
 326 };
 327
 328 static const AVFilterPad deint_cuda_inputs[] = {
 329     {
 330         .name          = "default",
 331         .type          = AVMEDIA_TYPE_VIDEO,
 332         .filter_frame  = ff_yadif_filter_frame,
 333         .config_props  = config_input,
 334     },
 335 };
 336
 337 static const AVFilterPad deint_cuda_outputs[] = {
 338     {
 339         .name          = "default",
 340         .type          = AVMEDIA_TYPE_VIDEO,
 341         .request_frame = ff_yadif_request_frame,
 342         .config_props  = config_output,
 343     },
 344 };
 345
 346 const FFFilter ff_vf_yadif_cuda = {
 347     .p.name         = "yadif_cuda",
 348     .p.description  = NULL_IF_CONFIG_SMALL("Deinterlace CUDA frames"),
 349     .p.priv_class   = &yadif_cuda_class,
 350     .p.flags        = AVFILTER_FLAG_SUPPORT_TIMELINE_INTERNAL,
 351     .priv_size      = sizeof(DeintCUDAContext),
 352     .uninit         = deint_cuda_uninit,
 353     FILTER_SINGLE_PIXFMT(AV_PIX_FMT_CUDA),
 354     FILTER_INPUTS(deint_cuda_inputs),
 355     FILTER_OUTPUTS(deint_cuda_outputs),
 356     .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
 357 };