FFmpeg
vf_colorspace_cuda.c
Go to the documentation of this file.
1 /*
2  * Copyright (c) 2022, 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 <string.h>
24 
25 #include "libavutil/avstring.h"
26 #include "libavutil/common.h"
27 #include "libavutil/cuda_check.h"
28 #include "libavutil/hwcontext.h"
30 #include "libavutil/internal.h"
31 #include "libavutil/opt.h"
32 #include "libavutil/pixdesc.h"
33 
34 #include "avfilter.h"
35 #include "formats.h"
36 #include "internal.h"
37 #include "scale_eval.h"
38 #include "video.h"
39 
40 #include "cuda/load_helper.h"
41 
42 static const enum AVPixelFormat supported_formats[] = {
46 };
47 
48 #define DIV_UP(a, b) (((a) + (b)-1) / (b))
49 #define BLOCKX 32
50 #define BLOCKY 16
51 
52 #define CHECK_CU(x) FF_CUDA_CHECK_DL(ctx, s->hwctx->internal->cuda_dl, x)
53 
54 typedef struct CUDAColorspaceContext {
55  const AVClass* class;
56 
61 
62  CUcontext cu_ctx;
63  CUstream cu_stream;
64  CUmodule cu_module;
66 
69 
72 
74 {
75  CUDAColorspaceContext* s = ctx->priv;
76 
77  s->own_frame = av_frame_alloc();
78  if (!s->own_frame)
79  return AVERROR(ENOMEM);
80 
81  s->tmp_frame = av_frame_alloc();
82  if (!s->tmp_frame)
83  return AVERROR(ENOMEM);
84 
85  return 0;
86 }
87 
89 {
90  CUDAColorspaceContext* s = ctx->priv;
91 
92  if (s->hwctx && s->cu_module) {
93  CudaFunctions* cu = s->hwctx->internal->cuda_dl;
94  CUcontext dummy;
95 
96  CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx));
97  CHECK_CU(cu->cuModuleUnload(s->cu_module));
98  s->cu_module = NULL;
99  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
100  }
101 
102  av_frame_free(&s->own_frame);
103  av_buffer_unref(&s->frames_ctx);
104  av_frame_free(&s->tmp_frame);
105 }
106 
108  int width, int height)
109 {
110  AVBufferRef* out_ref = NULL;
111  AVHWFramesContext* out_ctx;
112  int ret;
113 
114  out_ref = av_hwframe_ctx_alloc(device_ctx);
115  if (!out_ref)
116  return AVERROR(ENOMEM);
117 
118  out_ctx = (AVHWFramesContext*)out_ref->data;
119 
120  out_ctx->format = AV_PIX_FMT_CUDA;
121  out_ctx->sw_format = s->pix_fmt;
122  out_ctx->width = FFALIGN(width, 32);
123  out_ctx->height = FFALIGN(height, 32);
124 
125  ret = av_hwframe_ctx_init(out_ref);
126  if (ret < 0)
127  goto fail;
128 
129  av_frame_unref(s->own_frame);
130  ret = av_hwframe_get_buffer(out_ref, s->own_frame, 0);
131  if (ret < 0)
132  goto fail;
133 
134  s->own_frame->width = width;
135  s->own_frame->height = height;
136 
137  av_buffer_unref(&s->frames_ctx);
138  s->frames_ctx = out_ref;
139 
140  return 0;
141 fail:
142  av_buffer_unref(&out_ref);
143  return ret;
144 }
145 
146 static int format_is_supported(enum AVPixelFormat fmt)
147 {
148  for (int i = 0; i < FF_ARRAY_ELEMS(supported_formats); i++)
149  if (fmt == supported_formats[i])
150  return 1;
151 
152  return 0;
153 }
154 
156  int height)
157 {
158  CUDAColorspaceContext* s = ctx->priv;
159  AVHWFramesContext* in_frames_ctx;
160 
161  int ret;
162 
163  if (!ctx->inputs[0]->hw_frames_ctx) {
164  av_log(ctx, AV_LOG_ERROR, "No hw context provided on input\n");
165  return AVERROR(EINVAL);
166  }
167 
168  in_frames_ctx = (AVHWFramesContext*)ctx->inputs[0]->hw_frames_ctx->data;
169  s->pix_fmt = in_frames_ctx->sw_format;
170 
171  if (!format_is_supported(s->pix_fmt)) {
172  av_log(ctx, AV_LOG_ERROR, "Unsupported pixel format: %s\n",
173  av_get_pix_fmt_name(s->pix_fmt));
174  return AVERROR(EINVAL);
175  }
176 
177  if ((AVCOL_RANGE_MPEG != s->range) && (AVCOL_RANGE_JPEG != s->range)) {
178  av_log(ctx, AV_LOG_ERROR, "Unsupported color range\n");
179  return AVERROR(EINVAL);
180  }
181 
182  s->num_planes = av_pix_fmt_count_planes(s->pix_fmt);
183 
184  ret = init_hwframe_ctx(s, in_frames_ctx->device_ref, width, height);
185  if (ret < 0)
186  return ret;
187 
188  ctx->outputs[0]->hw_frames_ctx = av_buffer_ref(s->frames_ctx);
189  if (!ctx->outputs[0]->hw_frames_ctx)
190  return AVERROR(ENOMEM);
191 
192  return 0;
193 }
194 
196 {
197  CUDAColorspaceContext* s = ctx->priv;
198  CUcontext dummy, cuda_ctx = s->hwctx->cuda_ctx;
199  CudaFunctions* cu = s->hwctx->internal->cuda_dl;
200  int ret;
201 
202  extern const unsigned char ff_vf_colorspace_cuda_ptx_data[];
203  extern const unsigned int ff_vf_colorspace_cuda_ptx_len;
204 
205  ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx));
206  if (ret < 0)
207  return ret;
208 
209  ret = ff_cuda_load_module(ctx, s->hwctx, &s->cu_module,
210  ff_vf_colorspace_cuda_ptx_data,
211  ff_vf_colorspace_cuda_ptx_len);
212  if (ret < 0)
213  goto fail;
214 
215  ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_convert[AVCOL_RANGE_MPEG], s->cu_module, "to_mpeg_cuda"));
216  if (ret < 0)
217  goto fail;
218 
219  ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_convert[AVCOL_RANGE_JPEG], s->cu_module, "to_jpeg_cuda"));
220  if (ret < 0)
221  goto fail;
222 
223 fail:
224  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
225  return ret;
226 }
227 
229 {
230  AVFilterContext* ctx = outlink->src;
231  AVFilterLink* inlink = outlink->src->inputs[0];
232  CUDAColorspaceContext* s = ctx->priv;
233  AVHWFramesContext* frames_ctx =
234  (AVHWFramesContext*)inlink->hw_frames_ctx->data;
235  AVCUDADeviceContext* device_hwctx = frames_ctx->device_ctx->hwctx;
236  int ret;
237 
238  s->hwctx = device_hwctx;
239  s->cu_stream = s->hwctx->stream;
240 
241  outlink->w = inlink->w;
242  outlink->h = inlink->h;
243 
245  if (ret < 0)
246  return ret;
247 
248  if (inlink->sample_aspect_ratio.num) {
249  outlink->sample_aspect_ratio = av_mul_q(
250  (AVRational){outlink->h * inlink->w, outlink->w * inlink->h},
251  inlink->sample_aspect_ratio);
252  } else {
253  outlink->sample_aspect_ratio = inlink->sample_aspect_ratio;
254  }
255 
257  if (ret < 0)
258  return ret;
259 
260  return ret;
261 }
262 
264 {
265  CUDAColorspaceContext* s = ctx->priv;
266  CudaFunctions* cu = s->hwctx->internal->cuda_dl;
267  CUcontext dummy, cuda_ctx = s->hwctx->cuda_ctx;
268  int ret;
269 
270  ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx));
271  if (ret < 0)
272  return ret;
273 
274  out->color_range = s->range;
275 
276  for (int i = 0; i < s->num_planes; i++) {
277  int width = in->width, height = in->height, comp_id = (i > 0);
278 
279  switch (s->pix_fmt) {
280  case AV_PIX_FMT_YUV444P:
281  break;
282  case AV_PIX_FMT_YUV420P:
283  width = comp_id ? in->width / 2 : in->width;
284  /* fall-through */
285  case AV_PIX_FMT_NV12:
286  height = comp_id ? in->height / 2 : in->height;
287  break;
288  default:
289  av_log(ctx, AV_LOG_ERROR, "Unsupported pixel format: %s\n",
290  av_get_pix_fmt_name(s->pix_fmt));
291  return AVERROR(EINVAL);
292  }
293 
294  if (!s->cu_convert[out->color_range]) {
295  av_log(ctx, AV_LOG_ERROR, "Unsupported color range\n");
296  return AVERROR(EINVAL);
297  }
298 
299  if (in->color_range != out->color_range) {
300  void* args[] = {&in->data[i], &out->data[i], &in->linesize[i],
301  &comp_id};
302  ret = CHECK_CU(cu->cuLaunchKernel(
303  s->cu_convert[out->color_range], DIV_UP(width, BLOCKX),
304  DIV_UP(height, BLOCKY), 1, BLOCKX, BLOCKY, 1, 0, s->cu_stream,
305  args, NULL));
306  } else {
307  ret = av_hwframe_transfer_data(out, in, 0);
308  if (ret < 0)
309  return ret;
310  }
311  }
312 
313  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
314  return ret;
315 }
316 
318 {
319  CUDAColorspaceContext* s = ctx->priv;
320  AVFilterLink* outlink = ctx->outputs[0];
321  AVFrame* src = in;
322  int ret;
323 
324  ret = conv_cuda_convert(ctx, s->own_frame, src);
325  if (ret < 0)
326  return ret;
327 
328  src = s->own_frame;
329  ret = av_hwframe_get_buffer(src->hw_frames_ctx, s->tmp_frame, 0);
330  if (ret < 0)
331  return ret;
332 
333  av_frame_move_ref(out, s->own_frame);
334  av_frame_move_ref(s->own_frame, s->tmp_frame);
335 
336  s->own_frame->width = outlink->w;
337  s->own_frame->height = outlink->h;
338 
339  ret = av_frame_copy_props(out, in);
340  if (ret < 0)
341  return ret;
342 
343  return 0;
344 }
345 
347 {
348  AVFilterContext* ctx = link->dst;
349  CUDAColorspaceContext* s = ctx->priv;
350  AVFilterLink* outlink = ctx->outputs[0];
351  CudaFunctions* cu = s->hwctx->internal->cuda_dl;
352 
353  AVFrame* out = NULL;
354  CUcontext dummy;
355  int ret = 0;
356 
357  out = av_frame_alloc();
358  if (!out) {
359  ret = AVERROR(ENOMEM);
360  goto fail;
361  }
362 
363  ret = CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx));
364  if (ret < 0)
365  goto fail;
366 
367  ret = cudacolorspace_conv(ctx, out, in);
368 
369  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
370  if (ret < 0)
371  goto fail;
372 
373  av_reduce(&out->sample_aspect_ratio.num, &out->sample_aspect_ratio.den,
374  (int64_t)in->sample_aspect_ratio.num * outlink->h * link->w,
375  (int64_t)in->sample_aspect_ratio.den * outlink->w * link->h,
376  INT_MAX);
377 
378  av_frame_free(&in);
379  return ff_filter_frame(outlink, out);
380 fail:
381  av_frame_free(&in);
382  av_frame_free(&out);
383  return ret;
384 }
385 
386 #define OFFSET(x) offsetof(CUDAColorspaceContext, x)
387 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
388 static const AVOption options[] = {
389  {"range", "Output video range", OFFSET(range), AV_OPT_TYPE_INT, { .i64 = AVCOL_RANGE_UNSPECIFIED }, AVCOL_RANGE_UNSPECIFIED, AVCOL_RANGE_NB - 1, FLAGS, "range"},
390  {"tv", "Limited range", 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_RANGE_MPEG }, 0, 0, FLAGS, "range"},
391  {"mpeg", "Limited range", 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_RANGE_MPEG }, 0, 0, FLAGS, "range"},
392  {"pc", "Full range", 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_RANGE_JPEG }, 0, 0, FLAGS, "range"},
393  {"jpeg", "Full range", 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_RANGE_JPEG }, 0, 0, FLAGS, "range"},
394  {NULL},
395 };
396 
398  .class_name = "colorspace_cuda",
399  .item_name = av_default_item_name,
400  .option = options,
401  .version = LIBAVUTIL_VERSION_INT,
402 };
403 
405  {
406  .name = "default",
407  .type = AVMEDIA_TYPE_VIDEO,
408  .filter_frame = cudacolorspace_filter_frame,
409  },
410 };
411 
413  {
414  .name = "default",
415  .type = AVMEDIA_TYPE_VIDEO,
416  .config_props = cudacolorspace_config_props,
417  },
418 };
419 
421  .name = "colorspace_cuda",
422  .description = NULL_IF_CONFIG_SMALL("CUDA accelerated video color converter"),
423 
424  .init = cudacolorspace_init,
425  .uninit = cudacolorspace_uninit,
426 
427  .priv_size = sizeof(CUDAColorspaceContext),
428  .priv_class = &cudacolorspace_class,
429 
432 
434 
435  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
436 };
AVHWDeviceContext::hwctx
void * hwctx
The format-specific data, allocated and freed by libavutil along with this context.
Definition: hwcontext.h:92
AVFrame::color_range
enum AVColorRange color_range
MPEG vs JPEG YUV range.
Definition: frame.h:578
AV_PIX_FMT_CUDA
@ AV_PIX_FMT_CUDA
HW acceleration through CUDA.
Definition: pixfmt.h:253
AVPixelFormat
AVPixelFormat
Pixel format.
Definition: pixfmt.h:64
AVERROR
Filter the word “frame” indicates either a video frame or a group of audio as stored in an AVFrame structure Format for each input and each output the list of supported formats For video that means pixel format For audio that means channel sample they are references to shared objects When the negotiation mechanism computes the intersection of the formats supported at each end of a all references to both lists are replaced with a reference to the intersection And when a single format is eventually chosen for a link amongst the remaining all references to the list are updated That means that if a filter requires that its input and output have the same format amongst a supported all it has to do is use a reference to the same list of formats query_formats can leave some formats unset and return AVERROR(EAGAIN) to cause the negotiation mechanism toagain later. That can be used by filters with complex requirements to use the format negotiated on one link to set the formats supported on another. Frame references ownership and permissions
opt.h
hwcontext_cuda_internal.h
out
FILE * out
Definition: movenc.c:54
FF_FILTER_FLAG_HWFRAME_AWARE
#define FF_FILTER_FLAG_HWFRAME_AWARE
The filter is aware of hardware frames, and any hardware frame context should not be automatically pr...
Definition: internal.h:370
ff_filter_frame
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition: avfilter.c:1009
AVBufferRef::data
uint8_t * data
The data buffer.
Definition: buffer.h:90
AVHWFramesContext::format
enum AVPixelFormat format
The pixel format identifying the underlying HW surface type.
Definition: hwcontext.h:209
cudacolorspace_config_props
static av_cold int cudacolorspace_config_props(AVFilterLink *outlink)
Definition: vf_colorspace_cuda.c:228
ff_cuda_load_module
int ff_cuda_load_module(void *avctx, AVCUDADeviceContext *hwctx, CUmodule *cu_module, const unsigned char *data, const unsigned int length)
Loads a CUDA module and applies any decompression, if neccesary.
Definition: load_helper.c:34
inlink
The exact code depends on how similar the blocks are and how related they are to the and needs to apply these operations to the correct inlink or outlink if there are several Macros are available to factor that when no extra processing is inlink
Definition: filter_design.txt:212
av_frame_free
void av_frame_free(AVFrame **frame)
Free the frame and any dynamically allocated objects in it, e.g.
Definition: frame.c:116
cudacolorspace_uninit
static av_cold void cudacolorspace_uninit(AVFilterContext *ctx)
Definition: vf_colorspace_cuda.c:88
av_hwframe_ctx_init
int av_hwframe_ctx_init(AVBufferRef *ref)
Finalize the context before use.
Definition: hwcontext.c:334
AVFrame
This structure describes decoded (raw) audio or video data.
Definition: frame.h:325
pixdesc.h
CUDAColorspaceContext::num_planes
int num_planes
Definition: vf_colorspace_cuda.c:70
AVFrame::width
int width
Definition: frame.h:397
AVCOL_RANGE_JPEG
@ AVCOL_RANGE_JPEG
Full range content.
Definition: pixfmt.h:661
CUDAColorspaceContext::tmp_frame
AVFrame * tmp_frame
Definition: vf_colorspace_cuda.c:60
av_hwframe_ctx_alloc
AVBufferRef * av_hwframe_ctx_alloc(AVBufferRef *device_ref_in)
Allocate an AVHWFramesContext tied to a given device context.
Definition: hwcontext.c:248
AVOption
AVOption.
Definition: opt.h:251
OFFSET
#define OFFSET(x)
Definition: vf_colorspace_cuda.c:386
av_buffer_ref
AVBufferRef * av_buffer_ref(const AVBufferRef *buf)
Create a new reference to an AVBuffer.
Definition: buffer.c:103
AVFilter::name
const char * name
Filter name.
Definition: avfilter.h:175
CUDAColorspaceContext::cu_convert
CUfunction cu_convert[AVCOL_RANGE_NB]
Definition: vf_colorspace_cuda.c:65
AVHWFramesContext::width
int width
The allocated dimensions of the frames in this pool.
Definition: hwcontext.h:229
video.h
CUDAColorspaceContext::pix_fmt
enum AVPixelFormat pix_fmt
Definition: vf_colorspace_cuda.c:67
DIV_UP
#define DIV_UP(a, b)
Definition: vf_colorspace_cuda.c:48
AVFrame::data
uint8_t * data[AV_NUM_DATA_POINTERS]
pointer to the picture/channel planes.
Definition: frame.h:346
FLAGS
#define FLAGS
Definition: vf_colorspace_cuda.c:387
formats.h
BLOCKY
#define BLOCKY
Definition: vf_colorspace_cuda.c:50
av_pix_fmt_count_planes
int av_pix_fmt_count_planes(enum AVPixelFormat pix_fmt)
Definition: pixdesc.c:2898
format_is_supported
static int format_is_supported(enum AVPixelFormat fmt)
Definition: vf_colorspace_cuda.c:146
fail
#define fail()
Definition: checkasm.h:134
dummy
int dummy
Definition: motion.c:65
AVCOL_RANGE_NB
@ AVCOL_RANGE_NB
Not part of ABI.
Definition: pixfmt.h:662
av_reduce
int av_reduce(int *dst_num, int *dst_den, int64_t num, int64_t den, int64_t max)
Reduce a fraction.
Definition: rational.c:35
AVRational::num
int num
Numerator.
Definition: rational.h:59
AVFilterPad
A filter pad used for either input or output.
Definition: internal.h:49
av_frame_alloc
AVFrame * av_frame_alloc(void)
Allocate an AVFrame and set its fields to default values.
Definition: frame.c:104
CUDAColorspaceContext::cu_module
CUmodule cu_module
Definition: vf_colorspace_cuda.c:64
AV_LOG_ERROR
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition: log.h:180
FF_ARRAY_ELEMS
#define FF_ARRAY_ELEMS(a)
Definition: sinewin_tablegen.c:29
av_cold
#define av_cold
Definition: attributes.h:90
AVHWFramesContext::height
int height
Definition: hwcontext.h:229
width
#define width
s
#define s(width, name)
Definition: cbs_vp9.c:256
CUDAColorspaceContext::own_frame
AVFrame * own_frame
Definition: vf_colorspace_cuda.c:59
ctx
AVFormatContext * ctx
Definition: movenc.c:48
load_helper.h
AV_PIX_FMT_YUV420P
@ AV_PIX_FMT_YUV420P
planar YUV 4:2:0, 12bpp, (1 Cr & Cb sample per 2x2 Y samples)
Definition: pixfmt.h:66
FILTER_INPUTS
#define FILTER_INPUTS(array)
Definition: internal.h:190
link
Filter the word “frame” indicates either a video frame or a group of audio as stored in an AVFrame structure Format for each input and each output the list of supported formats For video that means pixel format For audio that means channel sample they are references to shared objects When the negotiation mechanism computes the intersection of the formats supported at each end of a link
Definition: filter_design.txt:23
if
if(ret)
Definition: filter_design.txt:179
LIBAVUTIL_VERSION_INT
#define LIBAVUTIL_VERSION_INT
Definition: version.h:85
AVClass
Describe the class of an AVClass context structure.
Definition: log.h:66
NULL
#define NULL
Definition: coverity.c:32
AVHWFramesContext::sw_format
enum AVPixelFormat sw_format
The pixel format identifying the actual data layout of the hardware frames.
Definition: hwcontext.h:222
av_frame_copy_props
int av_frame_copy_props(AVFrame *dst, const AVFrame *src)
Copy only "metadata" fields from src to dst.
Definition: frame.c:603
av_buffer_unref
void av_buffer_unref(AVBufferRef **buf)
Free a given reference and automatically free the buffer if there are no more references to it.
Definition: buffer.c:139
CUDAColorspaceContext::cu_stream
CUstream cu_stream
Definition: vf_colorspace_cuda.c:63
AVRational
Rational number (pair of numerator and denominator).
Definition: rational.h:58
AVHWFramesContext::device_ref
AVBufferRef * device_ref
A reference to the parent AVHWDeviceContext.
Definition: hwcontext.h:141
AVFilterContext::inputs
AVFilterLink ** inputs
array of pointers to input links
Definition: avfilter.h:423
av_default_item_name
const char * av_default_item_name(void *ptr)
Return the context name.
Definition: log.c:237
CUDAColorspaceContext::cu_ctx
CUcontext cu_ctx
Definition: vf_colorspace_cuda.c:62
CUDAColorspaceContext
Definition: vf_colorspace_cuda.c:54
cudacolorspace_conv
static int cudacolorspace_conv(AVFilterContext *ctx, AVFrame *out, AVFrame *in)
Definition: vf_colorspace_cuda.c:317
AVCOL_RANGE_UNSPECIFIED
@ AVCOL_RANGE_UNSPECIFIED
Definition: pixfmt.h:627
NULL_IF_CONFIG_SMALL
#define NULL_IF_CONFIG_SMALL(x)
Return NULL if CONFIG_SMALL is true, otherwise the argument without modification.
Definition: internal.h:115
supported_formats
static enum AVPixelFormat supported_formats[]
Definition: vf_colorspace_cuda.c:42
scale_eval.h
BLOCKX
#define BLOCKX
Definition: vf_colorspace_cuda.c:49
height
#define height
internal.h
FILTER_SINGLE_PIXFMT
#define FILTER_SINGLE_PIXFMT(pix_fmt_)
Definition: internal.h:180
CHECK_CU
#define CHECK_CU(x)
Definition: vf_colorspace_cuda.c:52
i
#define i(width, name, range_min, range_max)
Definition: cbs_h2645.c:269
cudacolorspace_inputs
static const AVFilterPad cudacolorspace_inputs[]
Definition: vf_colorspace_cuda.c:404
internal.h
common.h
av_frame_move_ref
void av_frame_move_ref(AVFrame *dst, AVFrame *src)
Move everything contained in src to dst and reset src.
Definition: frame.c:516
av_frame_unref
void av_frame_unref(AVFrame *frame)
Unreference all the buffers referenced by frame and reset the frame fields.
Definition: frame.c:487
init_hwframe_ctx
static av_cold int init_hwframe_ctx(CUDAColorspaceContext *s, AVBufferRef *device_ctx, int width, int height)
Definition: vf_colorspace_cuda.c:107
AVFilterPad::name
const char * name
Pad name.
Definition: internal.h:55
AVCOL_RANGE_MPEG
@ AVCOL_RANGE_MPEG
Narrow or limited range content.
Definition: pixfmt.h:644
ff_vf_colorspace_cuda
const AVFilter ff_vf_colorspace_cuda
Definition: vf_colorspace_cuda.c:420
AVFilter
Filter definition.
Definition: avfilter.h:171
AVHWFramesContext
This struct describes a set or pool of "hardware" frames (i.e.
Definition: hwcontext.h:124
AVCUDADeviceContext
This struct is allocated as AVHWDeviceContext.hwctx.
Definition: hwcontext_cuda.h:42
ret
ret
Definition: filter_design.txt:187
AV_PIX_FMT_NV12
@ AV_PIX_FMT_NV12
planar YUV 4:2:0, 12bpp, 1 plane for Y and 1 plane for the UV components, which are interleaved (firs...
Definition: pixfmt.h:89
AVClass::class_name
const char * class_name
The name of the class; usually it is the same name as the context structure type to which the AVClass...
Definition: log.h:71
cudacolorspace_init
static av_cold int cudacolorspace_init(AVFilterContext *ctx)
Definition: vf_colorspace_cuda.c:73
AVHWFramesContext::device_ctx
AVHWDeviceContext * device_ctx
The parent AVHWDeviceContext.
Definition: hwcontext.h:149
cuda_check.h
AVFrame::sample_aspect_ratio
AVRational sample_aspect_ratio
Sample aspect ratio for the video frame, 0/1 if unknown/unspecified.
Definition: frame.h:427
av_hwframe_transfer_data
int av_hwframe_transfer_data(AVFrame *dst, const AVFrame *src, int flags)
Copy data to or from a hw surface.
Definition: hwcontext.c:448
AVFrame::height
int height
Definition: frame.h:397
init_processing_chain
static av_cold int init_processing_chain(AVFilterContext *ctx, int width, int height)
Definition: vf_colorspace_cuda.c:155
cudacolorspace_filter_frame
static int cudacolorspace_filter_frame(AVFilterLink *link, AVFrame *in)
Definition: vf_colorspace_cuda.c:346
AVRational::den
int den
Denominator.
Definition: rational.h:60
AV_OPT_TYPE_INT
@ AV_OPT_TYPE_INT
Definition: opt.h:225
avfilter.h
CUDAColorspaceContext::range
enum AVColorRange range
Definition: vf_colorspace_cuda.c:68
av_mul_q
AVRational av_mul_q(AVRational b, AVRational c)
Multiply two rationals.
Definition: rational.c:80
AV_PIX_FMT_YUV444P
@ AV_PIX_FMT_YUV444P
planar YUV 4:4:4, 24bpp, (1 Cr & Cb sample per 1x1 Y samples)
Definition: pixfmt.h:71
AVFilterContext
An instance of a filter.
Definition: avfilter.h:415
AVMEDIA_TYPE_VIDEO
@ AVMEDIA_TYPE_VIDEO
Definition: avutil.h:201
AVBufferRef
A reference to a data buffer.
Definition: buffer.h:82
conv_cuda_convert
static int conv_cuda_convert(AVFilterContext *ctx, AVFrame *out, AVFrame *in)
Definition: vf_colorspace_cuda.c:263
FFALIGN
#define FFALIGN(x, a)
Definition: macros.h:78
FILTER_OUTPUTS
#define FILTER_OUTPUTS(array)
Definition: internal.h:191
src
INIT_CLIP pixel * src
Definition: h264pred_template.c:418
CUDAColorspaceContext::frames_ctx
AVBufferRef * frames_ctx
Definition: vf_colorspace_cuda.c:58
hwcontext.h
AVFrame::linesize
int linesize[AV_NUM_DATA_POINTERS]
For video, a positive or negative value, which is typically indicating the size in bytes of each pict...
Definition: frame.h:370
av_log
#define av_log(a,...)
Definition: tableprint_vlc.h:27
cudacolorspace_class
static const AVClass cudacolorspace_class
Definition: vf_colorspace_cuda.c:397
avstring.h
cudacolorspace_outputs
static const AVFilterPad cudacolorspace_outputs[]
Definition: vf_colorspace_cuda.c:412
cudacolorspace_load_functions
static av_cold int cudacolorspace_load_functions(AVFilterContext *ctx)
Definition: vf_colorspace_cuda.c:195
AVColorRange
AVColorRange
Visual content value range.
Definition: pixfmt.h:626
av_hwframe_get_buffer
int av_hwframe_get_buffer(AVBufferRef *hwframe_ref, AVFrame *frame, int flags)
Allocate a new frame attached to the given AVHWFramesContext.
Definition: hwcontext.c:507
AV_OPT_TYPE_CONST
@ AV_OPT_TYPE_CONST
Definition: opt.h:234
CUDAColorspaceContext::hwctx
AVCUDADeviceContext * hwctx
Definition: vf_colorspace_cuda.c:57
av_get_pix_fmt_name
const char * av_get_pix_fmt_name(enum AVPixelFormat pix_fmt)
Return the short name for a pixel format, NULL in case pix_fmt is unknown.
Definition: pixdesc.c:2778
options
static const AVOption options[]
Definition: vf_colorspace_cuda.c:388