1 /*
2 * Copyright (c) 2017, NVIDIA CORPORATION. All rights reserved.
3 *
4 * Permission is hereby granted, free of charge, to any person obtaining a
5 * copy of this software and associated documentation files (the "Software"),
6 * to deal in the Software without restriction, including without limitation
7 * the rights to use, copy, modify, merge, publish, distribute, sublicense,
8 * and/or sell copies of the Software, and to permit persons to whom the
9 * Software is furnished to do so, subject to the following conditions:
10 *
11 * The above copyright notice and this permission notice shall be included in
12 * all copies or substantial portions of the Software.
13 *
14 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
15 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
16 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
17 * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
18 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
19 * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER
20 * DEALINGS IN THE SOFTWARE.
21 */
22
23 #include <float.h>
24 #include <stdio.h>
25 #include <string.h>
26
27 #include "libavutil/avstring.h"
28 #include "libavutil/common.h"
29 #include "libavutil/hwcontext.h"
30 #include "libavutil/hwcontext_cuda_internal.h"
31 #include "libavutil/cuda_check.h"
32 #include "libavutil/internal.h"
33 #include "libavutil/opt.h"
34 #include "libavutil/pixdesc.h"
35
36 #include "avfilter.h"
37 #include "formats.h"
38 #include "internal.h"
39 #include "scale_eval.h"
40 #include "video.h"
41
42 #include "cuda/load_helper.h"
43 #include "vf_scale_cuda.h"
44
45 static const enum AVPixelFormat supported_formats[] = {
46 AV_PIX_FMT_YUV420P,
47 AV_PIX_FMT_NV12,
48 AV_PIX_FMT_YUV444P,
49 AV_PIX_FMT_P010,
50 AV_PIX_FMT_P016,
51 AV_PIX_FMT_YUV444P16,
52 AV_PIX_FMT_0RGB32,
53 AV_PIX_FMT_0BGR32,
54 };
55
56 #define DIV_UP(a, b) ( ((a) + (b) - 1) / (b) )
57 #define BLOCKX 32
58 #define BLOCKY 16
59
60 #define CHECK_CU(x) FF_CUDA_CHECK_DL(ctx, s->hwctx->internal->cuda_dl, x)
61
62 enum {
63 INTERP_ALGO_DEFAULT,
64
65 INTERP_ALGO_NEAREST,
66 INTERP_ALGO_BILINEAR,
67 INTERP_ALGO_BICUBIC,
68 INTERP_ALGO_LANCZOS,
69
70 INTERP_ALGO_COUNT
71 };
72
73 typedef struct CUDAScaleContext {
74 const AVClass *class;
75
76 AVCUDADeviceContext *hwctx;
77
78 enum AVPixelFormat in_fmt, out_fmt;
79 const AVPixFmtDescriptor *in_desc, *out_desc;
80 int in_planes, out_planes;
81 int in_plane_depths[4];
82 int in_plane_channels[4];
83
84 AVBufferRef *frames_ctx;
85 AVFrame *frame;
86
87 AVFrame *tmp_frame;
88 int passthrough;
89
90 /**
91 * Output sw format. AV_PIX_FMT_NONE for no conversion.
92 */
93 enum AVPixelFormat format;
94
95 char *w_expr; ///< width expression string
96 char *h_expr; ///< height expression string
97
98 int force_original_aspect_ratio;
99 int force_divisible_by;
100
101 CUcontext cu_ctx;
102 CUmodule cu_module;
103 CUfunction cu_func;
104 CUfunction cu_func_uv;
105 CUstream cu_stream;
106
107 int interp_algo;
108 int interp_use_linear;
109 int interp_as_integer;
110
111 float param;
112 } CUDAScaleContext;
113
cudascale_init(AVFilterContext * ctx)114 static av_cold int cudascale_init(AVFilterContext *ctx)
115 {
116 CUDAScaleContext *s = ctx->priv;
117
118 s->frame = av_frame_alloc();
119 if (!s->frame)
120 return AVERROR(ENOMEM);
121
122 s->tmp_frame = av_frame_alloc();
123 if (!s->tmp_frame)
124 return AVERROR(ENOMEM);
125
126 return 0;
127 }
128
cudascale_uninit(AVFilterContext * ctx)129 static av_cold void cudascale_uninit(AVFilterContext *ctx)
130 {
131 CUDAScaleContext *s = ctx->priv;
132
133 if (s->hwctx && s->cu_module) {
134 CudaFunctions *cu = s->hwctx->internal->cuda_dl;
135 CUcontext dummy;
136
137 CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx));
138 CHECK_CU(cu->cuModuleUnload(s->cu_module));
139 s->cu_module = NULL;
140 CHECK_CU(cu->cuCtxPopCurrent(&dummy));
141 }
142
143 av_frame_free(&s->frame);
144 av_buffer_unref(&s->frames_ctx);
145 av_frame_free(&s->tmp_frame);
146 }
147
init_hwframe_ctx(CUDAScaleContext * s,AVBufferRef * device_ctx,int width,int height)148 static av_cold int init_hwframe_ctx(CUDAScaleContext *s, AVBufferRef *device_ctx, int width, int height)
149 {
150 AVBufferRef *out_ref = NULL;
151 AVHWFramesContext *out_ctx;
152 int ret;
153
154 out_ref = av_hwframe_ctx_alloc(device_ctx);
155 if (!out_ref)
156 return AVERROR(ENOMEM);
157 out_ctx = (AVHWFramesContext*)out_ref->data;
158
159 out_ctx->format = AV_PIX_FMT_CUDA;
160 out_ctx->sw_format = s->out_fmt;
161 out_ctx->width = FFALIGN(width, 32);
162 out_ctx->height = FFALIGN(height, 32);
163
164 ret = av_hwframe_ctx_init(out_ref);
165 if (ret < 0)
166 goto fail;
167
168 av_frame_unref(s->frame);
169 ret = av_hwframe_get_buffer(out_ref, s->frame, 0);
170 if (ret < 0)
171 goto fail;
172
173 s->frame->width = width;
174 s->frame->height = height;
175
176 av_buffer_unref(&s->frames_ctx);
177 s->frames_ctx = out_ref;
178
179 return 0;
180 fail:
181 av_buffer_unref(&out_ref);
182 return ret;
183 }
184
format_is_supported(enum AVPixelFormat fmt)185 static int format_is_supported(enum AVPixelFormat fmt)
186 {
187 int i;
188
189 for (i = 0; i < FF_ARRAY_ELEMS(supported_formats); i++)
190 if (supported_formats[i] == fmt)
191 return 1;
192 return 0;
193 }
194
set_format_info(AVFilterContext * ctx,enum AVPixelFormat in_format,enum AVPixelFormat out_format)195 static av_cold void set_format_info(AVFilterContext *ctx, enum AVPixelFormat in_format, enum AVPixelFormat out_format)
196 {
197 CUDAScaleContext *s = ctx->priv;
198 int i, p, d;
199
200 s->in_fmt = in_format;
201 s->out_fmt = out_format;
202
203 s->in_desc = av_pix_fmt_desc_get(s->in_fmt);
204 s->out_desc = av_pix_fmt_desc_get(s->out_fmt);
205 s->in_planes = av_pix_fmt_count_planes(s->in_fmt);
206 s->out_planes = av_pix_fmt_count_planes(s->out_fmt);
207
208 // find maximum step of each component of each plane
209 // For our subset of formats, this should accurately tell us how many channels CUDA needs
210 // i.e. 1 for Y plane, 2 for UV plane of NV12, 4 for single plane of RGB0 formats
211
212 for (i = 0; i < s->in_desc->nb_components; i++) {
213 d = (s->in_desc->comp[i].depth + 7) / 8;
214 p = s->in_desc->comp[i].plane;
215 s->in_plane_channels[p] = FFMAX(s->in_plane_channels[p], s->in_desc->comp[i].step / d);
216
217 s->in_plane_depths[p] = s->in_desc->comp[i].depth;
218 }
219 }
220
init_processing_chain(AVFilterContext * ctx,int in_width,int in_height,int out_width,int out_height)221 static av_cold int init_processing_chain(AVFilterContext *ctx, int in_width, int in_height,
222 int out_width, int out_height)
223 {
224 CUDAScaleContext *s = ctx->priv;
225
226 AVHWFramesContext *in_frames_ctx;
227
228 enum AVPixelFormat in_format;
229 enum AVPixelFormat out_format;
230 int ret;
231
232 /* check that we have a hw context */
233 if (!ctx->inputs[0]->hw_frames_ctx) {
234 av_log(ctx, AV_LOG_ERROR, "No hw context provided on input\n");
235 return AVERROR(EINVAL);
236 }
237 in_frames_ctx = (AVHWFramesContext*)ctx->inputs[0]->hw_frames_ctx->data;
238 in_format = in_frames_ctx->sw_format;
239 out_format = (s->format == AV_PIX_FMT_NONE) ? in_format : s->format;
240
241 if (!format_is_supported(in_format)) {
242 av_log(ctx, AV_LOG_ERROR, "Unsupported input format: %s\n",
243 av_get_pix_fmt_name(in_format));
244 return AVERROR(ENOSYS);
245 }
246 if (!format_is_supported(out_format)) {
247 av_log(ctx, AV_LOG_ERROR, "Unsupported output format: %s\n",
248 av_get_pix_fmt_name(out_format));
249 return AVERROR(ENOSYS);
250 }
251
252 set_format_info(ctx, in_format, out_format);
253
254 if (s->passthrough && in_width == out_width && in_height == out_height && in_format == out_format) {
255 s->frames_ctx = av_buffer_ref(ctx->inputs[0]->hw_frames_ctx);
256 if (!s->frames_ctx)
257 return AVERROR(ENOMEM);
258 } else {
259 s->passthrough = 0;
260
261 ret = init_hwframe_ctx(s, in_frames_ctx->device_ref, out_width, out_height);
262 if (ret < 0)
263 return ret;
264
265 if (in_width == out_width && in_height == out_height &&
266 in_format == out_format && s->interp_algo == INTERP_ALGO_DEFAULT)
267 s->interp_algo = INTERP_ALGO_NEAREST;
268 }
269
270 ctx->outputs[0]->hw_frames_ctx = av_buffer_ref(s->frames_ctx);
271 if (!ctx->outputs[0]->hw_frames_ctx)
272 return AVERROR(ENOMEM);
273
274 return 0;
275 }
276
cudascale_load_functions(AVFilterContext * ctx)277 static av_cold int cudascale_load_functions(AVFilterContext *ctx)
278 {
279 CUDAScaleContext *s = ctx->priv;
280 CUcontext dummy, cuda_ctx = s->hwctx->cuda_ctx;
281 CudaFunctions *cu = s->hwctx->internal->cuda_dl;
282 char buf[128];
283 int ret;
284
285 const char *in_fmt_name = av_get_pix_fmt_name(s->in_fmt);
286 const char *out_fmt_name = av_get_pix_fmt_name(s->out_fmt);
287
288 const char *function_infix = "";
289
290 extern const unsigned char ff_vf_scale_cuda_ptx_data[];
291 extern const unsigned int ff_vf_scale_cuda_ptx_len;
292
293 switch(s->interp_algo) {
294 case INTERP_ALGO_NEAREST:
295 function_infix = "Nearest";
296 s->interp_use_linear = 0;
297 s->interp_as_integer = 1;
298 break;
299 case INTERP_ALGO_BILINEAR:
300 function_infix = "Bilinear";
301 s->interp_use_linear = 1;
302 s->interp_as_integer = 1;
303 break;
304 case INTERP_ALGO_DEFAULT:
305 case INTERP_ALGO_BICUBIC:
306 function_infix = "Bicubic";
307 s->interp_use_linear = 0;
308 s->interp_as_integer = 0;
309 break;
310 case INTERP_ALGO_LANCZOS:
311 function_infix = "Lanczos";
312 s->interp_use_linear = 0;
313 s->interp_as_integer = 0;
314 break;
315 default:
316 av_log(ctx, AV_LOG_ERROR, "Unknown interpolation algorithm\n");
317 return AVERROR_BUG;
318 }
319
320 ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx));
321 if (ret < 0)
322 return ret;
323
324 ret = ff_cuda_load_module(ctx, s->hwctx, &s->cu_module,
325 ff_vf_scale_cuda_ptx_data, ff_vf_scale_cuda_ptx_len);
326 if (ret < 0)
327 goto fail;
328
329 snprintf(buf, sizeof(buf), "Subsample_%s_%s_%s", function_infix, in_fmt_name, out_fmt_name);
330 ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func, s->cu_module, buf));
331 if (ret < 0) {
332 av_log(ctx, AV_LOG_FATAL, "Unsupported conversion: %s -> %s\n", in_fmt_name, out_fmt_name);
333 ret = AVERROR(ENOSYS);
334 goto fail;
335 }
336
337 snprintf(buf, sizeof(buf), "Subsample_%s_%s_%s_uv", function_infix, in_fmt_name, out_fmt_name);
338 ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uv, s->cu_module, buf));
339 if (ret < 0)
340 goto fail;
341
342 fail:
343 CHECK_CU(cu->cuCtxPopCurrent(&dummy));
344
345 return ret;
346 }
347
cudascale_config_props(AVFilterLink * outlink)348 static av_cold int cudascale_config_props(AVFilterLink *outlink)
349 {
350 AVFilterContext *ctx = outlink->src;
351 AVFilterLink *inlink = outlink->src->inputs[0];
352 CUDAScaleContext *s = ctx->priv;
353 AVHWFramesContext *frames_ctx = (AVHWFramesContext*)inlink->hw_frames_ctx->data;
354 AVCUDADeviceContext *device_hwctx = frames_ctx->device_ctx->hwctx;
355 int w, h;
356 int ret;
357
358 s->hwctx = device_hwctx;
359 s->cu_stream = s->hwctx->stream;
360
361 if ((ret = ff_scale_eval_dimensions(s,
362 s->w_expr, s->h_expr,
363 inlink, outlink,
364 &w, &h)) < 0)
365 goto fail;
366
367 ff_scale_adjust_dimensions(inlink, &w, &h,
368 s->force_original_aspect_ratio, s->force_divisible_by);
369
370 if (((int64_t)h * inlink->w) > INT_MAX ||
371 ((int64_t)w * inlink->h) > INT_MAX)
372 av_log(ctx, AV_LOG_ERROR, "Rescaled value for width or height is too big.\n");
373
374 outlink->w = w;
375 outlink->h = h;
376
377 ret = init_processing_chain(ctx, inlink->w, inlink->h, w, h);
378 if (ret < 0)
379 return ret;
380
381 if (inlink->sample_aspect_ratio.num) {
382 outlink->sample_aspect_ratio = av_mul_q((AVRational){outlink->h*inlink->w,
383 outlink->w*inlink->h},
384 inlink->sample_aspect_ratio);
385 } else {
386 outlink->sample_aspect_ratio = inlink->sample_aspect_ratio;
387 }
388
389 av_log(ctx, AV_LOG_VERBOSE, "w:%d h:%d fmt:%s -> w:%d h:%d fmt:%s%s\n",
390 inlink->w, inlink->h, av_get_pix_fmt_name(s->in_fmt),
391 outlink->w, outlink->h, av_get_pix_fmt_name(s->out_fmt),
392 s->passthrough ? " (passthrough)" : "");
393
394 ret = cudascale_load_functions(ctx);
395 if (ret < 0)
396 return ret;
397
398 return 0;
399
400 fail:
401 return ret;
402 }
403
call_resize_kernel(AVFilterContext * ctx,CUfunction func,CUtexObject src_tex[4],int src_width,int src_height,AVFrame * out_frame,int dst_width,int dst_height,int dst_pitch)404 static int call_resize_kernel(AVFilterContext *ctx, CUfunction func,
405 CUtexObject src_tex[4], int src_width, int src_height,
406 AVFrame *out_frame, int dst_width, int dst_height, int dst_pitch)
407 {
408 CUDAScaleContext *s = ctx->priv;
409 CudaFunctions *cu = s->hwctx->internal->cuda_dl;
410
411 CUdeviceptr dst_devptr[4] = {
412 (CUdeviceptr)out_frame->data[0], (CUdeviceptr)out_frame->data[1],
413 (CUdeviceptr)out_frame->data[2], (CUdeviceptr)out_frame->data[3]
414 };
415
416 void *args_uchar[] = {
417 &src_tex[0], &src_tex[1], &src_tex[2], &src_tex[3],
418 &dst_devptr[0], &dst_devptr[1], &dst_devptr[2], &dst_devptr[3],
419 &dst_width, &dst_height, &dst_pitch,
420 &src_width, &src_height, &s->param
421 };
422
423 return CHECK_CU(cu->cuLaunchKernel(func,
424 DIV_UP(dst_width, BLOCKX), DIV_UP(dst_height, BLOCKY), 1,
425 BLOCKX, BLOCKY, 1, 0, s->cu_stream, args_uchar, NULL));
426 }
427
scalecuda_resize(AVFilterContext * ctx,AVFrame * out,AVFrame * in)428 static int scalecuda_resize(AVFilterContext *ctx,
429 AVFrame *out, AVFrame *in)
430 {
431 CUDAScaleContext *s = ctx->priv;
432 CudaFunctions *cu = s->hwctx->internal->cuda_dl;
433 CUcontext dummy, cuda_ctx = s->hwctx->cuda_ctx;
434 int i, ret;
435
436 CUtexObject tex[4] = { 0, 0, 0, 0 };
437
438 ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx));
439 if (ret < 0)
440 return ret;
441
442 for (i = 0; i < s->in_planes; i++) {
443 CUDA_TEXTURE_DESC tex_desc = {
444 .filterMode = s->interp_use_linear ?
445 CU_TR_FILTER_MODE_LINEAR :
446 CU_TR_FILTER_MODE_POINT,
447 .flags = s->interp_as_integer ? CU_TRSF_READ_AS_INTEGER : 0,
448 };
449
450 CUDA_RESOURCE_DESC res_desc = {
451 .resType = CU_RESOURCE_TYPE_PITCH2D,
452 .res.pitch2D.format = s->in_plane_depths[i] <= 8 ?
453 CU_AD_FORMAT_UNSIGNED_INT8 :
454 CU_AD_FORMAT_UNSIGNED_INT16,
455 .res.pitch2D.numChannels = s->in_plane_channels[i],
456 .res.pitch2D.pitchInBytes = in->linesize[i],
457 .res.pitch2D.devPtr = (CUdeviceptr)in->data[i],
458 };
459
460 if (i == 1 || i == 2) {
461 res_desc.res.pitch2D.width = AV_CEIL_RSHIFT(in->width, s->in_desc->log2_chroma_w);
462 res_desc.res.pitch2D.height = AV_CEIL_RSHIFT(in->height, s->in_desc->log2_chroma_h);
463 } else {
464 res_desc.res.pitch2D.width = in->width;
465 res_desc.res.pitch2D.height = in->height;
466 }
467
468 ret = CHECK_CU(cu->cuTexObjectCreate(&tex[i], &res_desc, &tex_desc, NULL));
469 if (ret < 0)
470 goto exit;
471 }
472
473 // scale primary plane(s). Usually Y (and A), or single plane of RGB frames.
474 ret = call_resize_kernel(ctx, s->cu_func,
475 tex, in->width, in->height,
476 out, out->width, out->height, out->linesize[0]);
477 if (ret < 0)
478 goto exit;
479
480 if (s->out_planes > 1) {
481 // scale UV plane. Scale function sets both U and V plane, or singular interleaved plane.
482 ret = call_resize_kernel(ctx, s->cu_func_uv, tex,
483 AV_CEIL_RSHIFT(in->width, s->in_desc->log2_chroma_w),
484 AV_CEIL_RSHIFT(in->height, s->in_desc->log2_chroma_h),
485 out,
486 AV_CEIL_RSHIFT(out->width, s->out_desc->log2_chroma_w),
487 AV_CEIL_RSHIFT(out->height, s->out_desc->log2_chroma_h),
488 out->linesize[1]);
489 if (ret < 0)
490 goto exit;
491 }
492
493 exit:
494 for (i = 0; i < s->in_planes; i++)
495 if (tex[i])
496 CHECK_CU(cu->cuTexObjectDestroy(tex[i]));
497
498 CHECK_CU(cu->cuCtxPopCurrent(&dummy));
499
500 return ret;
501 }
502
cudascale_scale(AVFilterContext * ctx,AVFrame * out,AVFrame * in)503 static int cudascale_scale(AVFilterContext *ctx, AVFrame *out, AVFrame *in)
504 {
505 CUDAScaleContext *s = ctx->priv;
506 AVFilterLink *outlink = ctx->outputs[0];
507 AVFrame *src = in;
508 int ret;
509
510 ret = scalecuda_resize(ctx, s->frame, src);
511 if (ret < 0)
512 return ret;
513
514 src = s->frame;
515 ret = av_hwframe_get_buffer(src->hw_frames_ctx, s->tmp_frame, 0);
516 if (ret < 0)
517 return ret;
518
519 av_frame_move_ref(out, s->frame);
520 av_frame_move_ref(s->frame, s->tmp_frame);
521
522 s->frame->width = outlink->w;
523 s->frame->height = outlink->h;
524
525 ret = av_frame_copy_props(out, in);
526 if (ret < 0)
527 return ret;
528
529 return 0;
530 }
531
cudascale_filter_frame(AVFilterLink * link,AVFrame * in)532 static int cudascale_filter_frame(AVFilterLink *link, AVFrame *in)
533 {
534 AVFilterContext *ctx = link->dst;
535 CUDAScaleContext *s = ctx->priv;
536 AVFilterLink *outlink = ctx->outputs[0];
537 CudaFunctions *cu = s->hwctx->internal->cuda_dl;
538
539 AVFrame *out = NULL;
540 CUcontext dummy;
541 int ret = 0;
542
543 if (s->passthrough)
544 return ff_filter_frame(outlink, in);
545
546 out = av_frame_alloc();
547 if (!out) {
548 ret = AVERROR(ENOMEM);
549 goto fail;
550 }
551
552 ret = CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx));
553 if (ret < 0)
554 goto fail;
555
556 ret = cudascale_scale(ctx, out, in);
557
558 CHECK_CU(cu->cuCtxPopCurrent(&dummy));
559 if (ret < 0)
560 goto fail;
561
562 av_reduce(&out->sample_aspect_ratio.num, &out->sample_aspect_ratio.den,
563 (int64_t)in->sample_aspect_ratio.num * outlink->h * link->w,
564 (int64_t)in->sample_aspect_ratio.den * outlink->w * link->h,
565 INT_MAX);
566
567 av_frame_free(&in);
568 return ff_filter_frame(outlink, out);
569 fail:
570 av_frame_free(&in);
571 av_frame_free(&out);
572 return ret;
573 }
574
cudascale_get_video_buffer(AVFilterLink * inlink,int w,int h)575 static AVFrame *cudascale_get_video_buffer(AVFilterLink *inlink, int w, int h)
576 {
577 CUDAScaleContext *s = inlink->dst->priv;
578
579 return s->passthrough ?
580 ff_null_get_video_buffer (inlink, w, h) :
581 ff_default_get_video_buffer(inlink, w, h);
582 }
583
584 #define OFFSET(x) offsetof(CUDAScaleContext, x)
585 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM|AV_OPT_FLAG_VIDEO_PARAM)
586 static const AVOption options[] = {
587 { "w", "Output video width", OFFSET(w_expr), AV_OPT_TYPE_STRING, { .str = "iw" }, .flags = FLAGS },
588 { "h", "Output video height", OFFSET(h_expr), AV_OPT_TYPE_STRING, { .str = "ih" }, .flags = FLAGS },
589 { "interp_algo", "Interpolation algorithm used for resizing", OFFSET(interp_algo), AV_OPT_TYPE_INT, { .i64 = INTERP_ALGO_DEFAULT }, 0, INTERP_ALGO_COUNT - 1, FLAGS, "interp_algo" },
590 { "nearest", "nearest neighbour", 0, AV_OPT_TYPE_CONST, { .i64 = INTERP_ALGO_NEAREST }, 0, 0, FLAGS, "interp_algo" },
591 { "bilinear", "bilinear", 0, AV_OPT_TYPE_CONST, { .i64 = INTERP_ALGO_BILINEAR }, 0, 0, FLAGS, "interp_algo" },
592 { "bicubic", "bicubic", 0, AV_OPT_TYPE_CONST, { .i64 = INTERP_ALGO_BICUBIC }, 0, 0, FLAGS, "interp_algo" },
593 { "lanczos", "lanczos", 0, AV_OPT_TYPE_CONST, { .i64 = INTERP_ALGO_LANCZOS }, 0, 0, FLAGS, "interp_algo" },
594 { "format", "Output video pixel format", OFFSET(format), AV_OPT_TYPE_PIXEL_FMT, { .i64 = AV_PIX_FMT_NONE }, INT_MIN, INT_MAX, .flags=FLAGS },
595 { "passthrough", "Do not process frames at all if parameters match", OFFSET(passthrough), AV_OPT_TYPE_BOOL, { .i64 = 1 }, 0, 1, FLAGS },
596 { "param", "Algorithm-Specific parameter", OFFSET(param), AV_OPT_TYPE_FLOAT, { .dbl = SCALE_CUDA_PARAM_DEFAULT }, -FLT_MAX, FLT_MAX, FLAGS },
597 { "force_original_aspect_ratio", "decrease or increase w/h if necessary to keep the original AR", OFFSET(force_original_aspect_ratio), AV_OPT_TYPE_INT, { .i64 = 0 }, 0, 2, FLAGS, "force_oar" },
598 { "disable", NULL, 0, AV_OPT_TYPE_CONST, {.i64 = 0 }, 0, 0, FLAGS, "force_oar" },
599 { "decrease", NULL, 0, AV_OPT_TYPE_CONST, {.i64 = 1 }, 0, 0, FLAGS, "force_oar" },
600 { "increase", NULL, 0, AV_OPT_TYPE_CONST, {.i64 = 2 }, 0, 0, FLAGS, "force_oar" },
601 { "force_divisible_by", "enforce that the output resolution is divisible by a defined integer when force_original_aspect_ratio is used", OFFSET(force_divisible_by), AV_OPT_TYPE_INT, { .i64 = 1 }, 1, 256, FLAGS },
602 { NULL },
603 };
604
605 static const AVClass cudascale_class = {
606 .class_name = "cudascale",
607 .item_name = av_default_item_name,
608 .option = options,
609 .version = LIBAVUTIL_VERSION_INT,
610 };
611
612 static const AVFilterPad cudascale_inputs[] = {
613 {
614 .name = "default",
615 .type = AVMEDIA_TYPE_VIDEO,
616 .filter_frame = cudascale_filter_frame,
617 .get_buffer.video = cudascale_get_video_buffer,
618 },
619 };
620
621 static const AVFilterPad cudascale_outputs[] = {
622 {
623 .name = "default",
624 .type = AVMEDIA_TYPE_VIDEO,
625 .config_props = cudascale_config_props,
626 },
627 };
628
629 const AVFilter ff_vf_scale_cuda = {
630 .name = "scale_cuda",
631 .description = NULL_IF_CONFIG_SMALL("GPU accelerated video resizer"),
632
633 .init = cudascale_init,
634 .uninit = cudascale_uninit,
635
636 .priv_size = sizeof(CUDAScaleContext),
637 .priv_class = &cudascale_class,
638
639 FILTER_INPUTS(cudascale_inputs),
640 FILTER_OUTPUTS(cudascale_outputs),
641
642 FILTER_SINGLE_PIXFMT(AV_PIX_FMT_CUDA),
643
644 .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
645 };
646