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/common.h"
26 #include "libavutil/cuda_check.h"
27 #include "libavutil/hwcontext.h"
29 #include "libavutil/internal.h"
30 #include "libavutil/opt.h"
31 #include "libavutil/pixdesc.h"
32 
33 #include "avfilter.h"
34 #include "filters.h"
35 
36 #include "cuda/load_helper.h"
37 
38 static const enum AVPixelFormat supported_formats[] = {
42 };
43 
44 #define DIV_UP(a, b) (((a) + (b)-1) / (b))
45 #define BLOCKX 32
46 #define BLOCKY 16
47 
48 #define CHECK_CU(x) FF_CUDA_CHECK_DL(ctx, s->hwctx->internal->cuda_dl, x)
49 
50 typedef struct CUDAColorspaceContext {
51  const AVClass* class;
52 
57 
58  CUcontext cu_ctx;
59  CUstream cu_stream;
60  CUmodule cu_module;
62 
65 
68 
70 {
71  CUDAColorspaceContext* s = ctx->priv;
72 
73  s->own_frame = av_frame_alloc();
74  if (!s->own_frame)
75  return AVERROR(ENOMEM);
76 
77  s->tmp_frame = av_frame_alloc();
78  if (!s->tmp_frame)
79  return AVERROR(ENOMEM);
80 
81  return 0;
82 }
83 
85 {
86  CUDAColorspaceContext* s = ctx->priv;
87 
88  if (s->hwctx && s->cu_module) {
89  CudaFunctions* cu = s->hwctx->internal->cuda_dl;
90  CUcontext dummy;
91 
92  CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx));
93  CHECK_CU(cu->cuModuleUnload(s->cu_module));
94  s->cu_module = NULL;
95  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
96  }
97 
98  av_frame_free(&s->own_frame);
99  av_buffer_unref(&s->frames_ctx);
100  av_frame_free(&s->tmp_frame);
101 }
102 
104  int width, int height)
105 {
106  AVBufferRef* out_ref = NULL;
107  AVHWFramesContext* out_ctx;
108  int ret;
109 
110  out_ref = av_hwframe_ctx_alloc(device_ctx);
111  if (!out_ref)
112  return AVERROR(ENOMEM);
113 
114  out_ctx = (AVHWFramesContext*)out_ref->data;
115 
116  out_ctx->format = AV_PIX_FMT_CUDA;
117  out_ctx->sw_format = s->pix_fmt;
118  out_ctx->width = FFALIGN(width, 32);
119  out_ctx->height = FFALIGN(height, 32);
120 
121  ret = av_hwframe_ctx_init(out_ref);
122  if (ret < 0)
123  goto fail;
124 
125  av_frame_unref(s->own_frame);
126  ret = av_hwframe_get_buffer(out_ref, s->own_frame, 0);
127  if (ret < 0)
128  goto fail;
129 
130  s->own_frame->width = width;
131  s->own_frame->height = height;
132 
133  av_buffer_unref(&s->frames_ctx);
134  s->frames_ctx = out_ref;
135 
136  return 0;
137 fail:
138  av_buffer_unref(&out_ref);
139  return ret;
140 }
141 
142 static int format_is_supported(enum AVPixelFormat fmt)
143 {
144  for (int i = 0; i < FF_ARRAY_ELEMS(supported_formats); i++)
145  if (fmt == supported_formats[i])
146  return 1;
147 
148  return 0;
149 }
150 
152  int height)
153 {
154  FilterLink *inl = ff_filter_link(ctx->inputs[0]);
155  FilterLink *outl = ff_filter_link(ctx->outputs[0]);
156  CUDAColorspaceContext* s = ctx->priv;
157  AVHWFramesContext* in_frames_ctx;
158 
159  int ret;
160 
161  if (!inl->hw_frames_ctx) {
162  av_log(ctx, AV_LOG_ERROR, "No hw context provided on input\n");
163  return AVERROR(EINVAL);
164  }
165 
166  in_frames_ctx = (AVHWFramesContext*)inl->hw_frames_ctx->data;
167  s->pix_fmt = in_frames_ctx->sw_format;
168 
169  if (!format_is_supported(s->pix_fmt)) {
170  av_log(ctx, AV_LOG_ERROR, "Unsupported pixel format: %s\n",
171  av_get_pix_fmt_name(s->pix_fmt));
172  return AVERROR(EINVAL);
173  }
174 
175  if ((AVCOL_RANGE_MPEG != s->range) && (AVCOL_RANGE_JPEG != s->range)) {
176  av_log(ctx, AV_LOG_ERROR, "Unsupported color range\n");
177  return AVERROR(EINVAL);
178  }
179 
180  s->num_planes = av_pix_fmt_count_planes(s->pix_fmt);
181 
182  ret = init_hwframe_ctx(s, in_frames_ctx->device_ref, width, height);
183  if (ret < 0)
184  return ret;
185 
186  outl->hw_frames_ctx = av_buffer_ref(s->frames_ctx);
187  if (!outl->hw_frames_ctx)
188  return AVERROR(ENOMEM);
189 
190  return 0;
191 }
192 
194 {
195  CUDAColorspaceContext* s = ctx->priv;
196  CUcontext dummy, cuda_ctx = s->hwctx->cuda_ctx;
197  CudaFunctions* cu = s->hwctx->internal->cuda_dl;
198  int ret;
199 
200  extern const unsigned char ff_vf_colorspace_cuda_ptx_data[];
201  extern const unsigned int ff_vf_colorspace_cuda_ptx_len;
202 
203  ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx));
204  if (ret < 0)
205  return ret;
206 
207  ret = ff_cuda_load_module(ctx, s->hwctx, &s->cu_module,
208  ff_vf_colorspace_cuda_ptx_data,
209  ff_vf_colorspace_cuda_ptx_len);
210  if (ret < 0)
211  goto fail;
212 
213  ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_convert[AVCOL_RANGE_MPEG], s->cu_module, "to_mpeg_cuda"));
214  if (ret < 0)
215  goto fail;
216 
217  ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_convert[AVCOL_RANGE_JPEG], s->cu_module, "to_jpeg_cuda"));
218  if (ret < 0)
219  goto fail;
220 
221 fail:
222  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
223  return ret;
224 }
225 
227 {
228  AVFilterContext* ctx = outlink->src;
229  AVFilterLink* inlink = outlink->src->inputs[0];
231  CUDAColorspaceContext* s = ctx->priv;
232  AVHWFramesContext* frames_ctx;
233  AVCUDADeviceContext* device_hwctx;
234  int ret;
235 
236  outlink->w = inlink->w;
237  outlink->h = inlink->h;
238 
240  if (ret < 0)
241  return ret;
242 
243  frames_ctx = (AVHWFramesContext*)inl->hw_frames_ctx->data;
244  device_hwctx = frames_ctx->device_ctx->hwctx;
245 
246  s->hwctx = device_hwctx;
247  s->cu_stream = s->hwctx->stream;
248 
249  if (inlink->sample_aspect_ratio.num) {
250  outlink->sample_aspect_ratio = av_mul_q(
251  (AVRational){outlink->h * inlink->w, outlink->w * inlink->h},
252  inlink->sample_aspect_ratio);
253  } else {
254  outlink->sample_aspect_ratio = inlink->sample_aspect_ratio;
255  }
256 
258  if (ret < 0)
259  return ret;
260 
261  return ret;
262 }
263 
265 {
266  CUDAColorspaceContext* s = ctx->priv;
267  CudaFunctions* cu = s->hwctx->internal->cuda_dl;
268  CUcontext dummy, cuda_ctx = s->hwctx->cuda_ctx;
269  int ret;
270 
271  ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx));
272  if (ret < 0)
273  return ret;
274 
275  out->color_range = s->range;
276 
277  for (int i = 0; i < s->num_planes; i++) {
278  int width = in->width, height = in->height, comp_id = (i > 0);
279 
280  switch (s->pix_fmt) {
281  case AV_PIX_FMT_YUV444P:
282  break;
283  case AV_PIX_FMT_YUV420P:
284  width = comp_id ? in->width / 2 : in->width;
285  /* fall-through */
286  case AV_PIX_FMT_NV12:
287  height = comp_id ? in->height / 2 : in->height;
288  break;
289  default:
290  av_log(ctx, AV_LOG_ERROR, "Unsupported pixel format: %s\n",
291  av_get_pix_fmt_name(s->pix_fmt));
292  return AVERROR(EINVAL);
293  }
294 
295  if (!s->cu_convert[out->color_range]) {
296  av_log(ctx, AV_LOG_ERROR, "Unsupported color range\n");
297  return AVERROR(EINVAL);
298  }
299 
300  if (in->color_range != out->color_range) {
301  void* args[] = {&in->data[i], &out->data[i], &in->linesize[i],
302  &comp_id};
303  ret = CHECK_CU(cu->cuLaunchKernel(
304  s->cu_convert[out->color_range], DIV_UP(width, BLOCKX),
305  DIV_UP(height, BLOCKY), 1, BLOCKX, BLOCKY, 1, 0, s->cu_stream,
306  args, NULL));
307  } else {
308  ret = av_hwframe_transfer_data(out, in, 0);
309  if (ret < 0)
310  return ret;
311  }
312  }
313 
314  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
315  return ret;
316 }
317 
319 {
320  CUDAColorspaceContext* s = ctx->priv;
321  AVFilterLink* outlink = ctx->outputs[0];
322  AVFrame* src = in;
323  int ret;
324 
325  ret = conv_cuda_convert(ctx, s->own_frame, src);
326  if (ret < 0)
327  return ret;
328 
329  src = s->own_frame;
330  ret = av_hwframe_get_buffer(src->hw_frames_ctx, s->tmp_frame, 0);
331  if (ret < 0)
332  return ret;
333 
334  av_frame_move_ref(out, s->own_frame);
335  av_frame_move_ref(s->own_frame, s->tmp_frame);
336 
337  s->own_frame->width = outlink->w;
338  s->own_frame->height = outlink->h;
339 
340  ret = av_frame_copy_props(out, in);
341  if (ret < 0)
342  return ret;
343 
344  return 0;
345 }
346 
348 {
349  AVFilterContext* ctx = link->dst;
350  CUDAColorspaceContext* s = ctx->priv;
351  AVFilterLink* outlink = ctx->outputs[0];
352  CudaFunctions* cu = s->hwctx->internal->cuda_dl;
353 
354  AVFrame* out = NULL;
355  CUcontext dummy;
356  int ret = 0;
357 
358  out = av_frame_alloc();
359  if (!out) {
360  ret = AVERROR(ENOMEM);
361  goto fail;
362  }
363 
364  ret = CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx));
365  if (ret < 0)
366  goto fail;
367 
368  ret = cudacolorspace_conv(ctx, out, in);
369 
370  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
371  if (ret < 0)
372  goto fail;
373 
374  av_reduce(&out->sample_aspect_ratio.num, &out->sample_aspect_ratio.den,
375  (int64_t)in->sample_aspect_ratio.num * outlink->h * link->w,
376  (int64_t)in->sample_aspect_ratio.den * outlink->w * link->h,
377  INT_MAX);
378 
379  av_frame_free(&in);
380  return ff_filter_frame(outlink, out);
381 fail:
382  av_frame_free(&in);
383  av_frame_free(&out);
384  return ret;
385 }
386 
387 #define OFFSET(x) offsetof(CUDAColorspaceContext, x)
388 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
389 static const AVOption options[] = {
390  {"range", "Output video range", OFFSET(range), AV_OPT_TYPE_INT, { .i64 = AVCOL_RANGE_UNSPECIFIED }, AVCOL_RANGE_UNSPECIFIED, AVCOL_RANGE_NB - 1, FLAGS, .unit = "range"},
391  {"tv", "Limited range", 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_RANGE_MPEG }, 0, 0, FLAGS, .unit = "range"},
392  {"mpeg", "Limited range", 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_RANGE_MPEG }, 0, 0, FLAGS, .unit = "range"},
393  {"pc", "Full range", 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_RANGE_JPEG }, 0, 0, FLAGS, .unit = "range"},
394  {"jpeg", "Full range", 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_RANGE_JPEG }, 0, 0, FLAGS, .unit = "range"},
395  {NULL},
396 };
397 
399  .class_name = "colorspace_cuda",
400  .item_name = av_default_item_name,
401  .option = options,
402  .version = LIBAVUTIL_VERSION_INT,
403 };
404 
406  {
407  .name = "default",
408  .type = AVMEDIA_TYPE_VIDEO,
409  .filter_frame = cudacolorspace_filter_frame,
410  },
411 };
412 
414  {
415  .name = "default",
416  .type = AVMEDIA_TYPE_VIDEO,
417  .config_props = cudacolorspace_config_props,
418  },
419 };
420 
422  .name = "colorspace_cuda",
423  .description = NULL_IF_CONFIG_SMALL("CUDA accelerated video color converter"),
424 
425  .init = cudacolorspace_init,
426  .uninit = cudacolorspace_uninit,
427 
428  .priv_size = sizeof(CUDAColorspaceContext),
429  .priv_class = &cudacolorspace_class,
430 
433 
435 
436  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
437 };
AVHWDeviceContext::hwctx
void * hwctx
The format-specific data, allocated and freed by libavutil along with this context.
Definition: hwcontext.h:85
AVFrame::color_range
enum AVColorRange color_range
MPEG vs JPEG YUV range.
Definition: frame.h:668
AV_PIX_FMT_CUDA
@ AV_PIX_FMT_CUDA
HW acceleration through CUDA.
Definition: pixfmt.h:260
AVPixelFormat
AVPixelFormat
Pixel format.
Definition: pixfmt.h:71
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:55
ff_filter_frame
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition: avfilter.c:1062
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:197
cudacolorspace_config_props
static av_cold int cudacolorspace_config_props(AVFilterLink *outlink)
Definition: vf_colorspace_cuda.c:226
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 necessary.
Definition: load_helper.c:35
int64_t
long long int64_t
Definition: coverity.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:162
cudacolorspace_uninit
static av_cold void cudacolorspace_uninit(AVFilterContext *ctx)
Definition: vf_colorspace_cuda.c:84
av_hwframe_ctx_init
int av_hwframe_ctx_init(AVBufferRef *ref)
Finalize the context before use.
Definition: hwcontext.c:322
FILTER_INPUTS
#define FILTER_INPUTS(array)
Definition: filters.h:262
AVFrame
This structure describes decoded (raw) audio or video data.
Definition: frame.h:389
pixdesc.h
CUDAColorspaceContext::num_planes
int num_planes
Definition: vf_colorspace_cuda.c:66
AVFrame::width
int width
Definition: frame.h:461
AVCOL_RANGE_JPEG
@ AVCOL_RANGE_JPEG
Full range content.
Definition: pixfmt.h:709
CUDAColorspaceContext::tmp_frame
AVFrame * tmp_frame
Definition: vf_colorspace_cuda.c:56
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:429
OFFSET
#define OFFSET(x)
Definition: vf_colorspace_cuda.c:387
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:205
CUDAColorspaceContext::cu_convert
CUfunction cu_convert[AVCOL_RANGE_NB]
Definition: vf_colorspace_cuda.c:61
AVHWFramesContext::width
int width
The allocated dimensions of the frames in this pool.
Definition: hwcontext.h:217
CUDAColorspaceContext::pix_fmt
enum AVPixelFormat pix_fmt
Definition: vf_colorspace_cuda.c:63
DIV_UP
#define DIV_UP(a, b)
Definition: vf_colorspace_cuda.c:44
AVFrame::data
uint8_t * data[AV_NUM_DATA_POINTERS]
pointer to the picture/channel planes.
Definition: frame.h:410
FLAGS
#define FLAGS
Definition: vf_colorspace_cuda.c:388
BLOCKY
#define BLOCKY
Definition: vf_colorspace_cuda.c:46
av_pix_fmt_count_planes
int av_pix_fmt_count_planes(enum AVPixelFormat pix_fmt)
Definition: pixdesc.c:3141
format_is_supported
static int format_is_supported(enum AVPixelFormat fmt)
Definition: vf_colorspace_cuda.c:142
fail
#define fail()
Definition: checkasm.h:188
dummy
int dummy
Definition: motion.c:66
AVCOL_RANGE_NB
@ AVCOL_RANGE_NB
Not part of ABI.
Definition: pixfmt.h:710
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: filters.h:38
av_frame_alloc
AVFrame * av_frame_alloc(void)
Allocate an AVFrame and set its fields to default values.
Definition: frame.c:150
CUDAColorspaceContext::cu_module
CUmodule cu_module
Definition: vf_colorspace_cuda.c:60
AV_LOG_ERROR
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition: log.h:209
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:217
s
#define s(width, name)
Definition: cbs_vp9.c:198
CUDAColorspaceContext::own_frame
AVFrame * own_frame
Definition: vf_colorspace_cuda.c:55
filters.h
ctx
AVFormatContext * ctx
Definition: movenc.c:49
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:73
FILTER_OUTPUTS
#define FILTER_OUTPUTS(array)
Definition: filters.h:263
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:75
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:210
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:713
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:59
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:126
AVFilterContext::inputs
AVFilterLink ** inputs
array of pointers to input links
Definition: avfilter.h:465
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:58
CUDAColorspaceContext
Definition: vf_colorspace_cuda.c:50
cudacolorspace_conv
static int cudacolorspace_conv(AVFilterContext *ctx, AVFrame *out, AVFrame *in)
Definition: vf_colorspace_cuda.c:318
AVCOL_RANGE_UNSPECIFIED
@ AVCOL_RANGE_UNSPECIFIED
Definition: pixfmt.h:675
ff_filter_link
static FilterLink * ff_filter_link(AVFilterLink *link)
Definition: filters.h:197
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: filters.h:206
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:94
height
#define height
Definition: dsp.h:85
supported_formats
static enum AVPixelFormat supported_formats[]
Definition: vf_colorspace_cuda.c:38
range
enum AVColorRange range
Definition: mediacodec_wrapper.c:2464
BLOCKX
#define BLOCKX
Definition: vf_colorspace_cuda.c:45
CHECK_CU
#define CHECK_CU(x)
Definition: vf_colorspace_cuda.c:48
i
#define i(width, name, range_min, range_max)
Definition: cbs_h2645.c:256
cudacolorspace_inputs
static const AVFilterPad cudacolorspace_inputs[]
Definition: vf_colorspace_cuda.c:405
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:637
av_frame_unref
void av_frame_unref(AVFrame *frame)
Unreference all the buffers referenced by frame and reset the frame fields.
Definition: frame.c:610
init_hwframe_ctx
static av_cold int init_hwframe_ctx(CUDAColorspaceContext *s, AVBufferRef *device_ctx, int width, int height)
Definition: vf_colorspace_cuda.c:103
AVFilterPad::name
const char * name
Pad name.
Definition: filters.h:44
AVCOL_RANGE_MPEG
@ AVCOL_RANGE_MPEG
Narrow or limited range content.
Definition: pixfmt.h:692
ff_vf_colorspace_cuda
const AVFilter ff_vf_colorspace_cuda
Definition: vf_colorspace_cuda.c:421
AVFilter
Filter definition.
Definition: avfilter.h:201
AVHWFramesContext
This struct describes a set or pool of "hardware" frames (i.e.
Definition: hwcontext.h:115
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:96
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:80
cudacolorspace_init
static av_cold int cudacolorspace_init(AVFilterContext *ctx)
Definition: vf_colorspace_cuda.c:69
AVHWFramesContext::device_ctx
AVHWDeviceContext * device_ctx
The parent AVHWDeviceContext.
Definition: hwcontext.h:134
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:496
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:433
AVFrame::height
int height
Definition: frame.h:461
init_processing_chain
static av_cold int init_processing_chain(AVFilterContext *ctx, int width, int height)
Definition: vf_colorspace_cuda.c:151
cudacolorspace_filter_frame
static int cudacolorspace_filter_frame(AVFilterLink *link, AVFrame *in)
Definition: vf_colorspace_cuda.c:347
AVRational::den
int den
Denominator.
Definition: rational.h:60
AV_OPT_TYPE_INT
@ AV_OPT_TYPE_INT
Underlying C type is int.
Definition: opt.h:259
avfilter.h
CUDAColorspaceContext::range
enum AVColorRange range
Definition: vf_colorspace_cuda.c:64
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:78
AVFilterContext
An instance of a filter.
Definition: avfilter.h:457
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:264
FFALIGN
#define FFALIGN(x, a)
Definition: macros.h:78
CUDAColorspaceContext::frames_ctx
AVBufferRef * frames_ctx
Definition: vf_colorspace_cuda.c:54
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:434
av_log
#define av_log(a,...)
Definition: tableprint_vlc.h:27
cudacolorspace_class
static const AVClass cudacolorspace_class
Definition: vf_colorspace_cuda.c:398
cudacolorspace_outputs
static const AVFilterPad cudacolorspace_outputs[]
Definition: vf_colorspace_cuda.c:413
width
#define width
Definition: dsp.h:85
cudacolorspace_load_functions
static av_cold int cudacolorspace_load_functions(AVFilterContext *ctx)
Definition: vf_colorspace_cuda.c:193
AVColorRange
AVColorRange
Visual content value range.
Definition: pixfmt.h:674
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:491
AV_OPT_TYPE_CONST
@ AV_OPT_TYPE_CONST
Special option type for declaring named constants.
Definition: opt.h:299
FILTER_SINGLE_PIXFMT
#define FILTER_SINGLE_PIXFMT(pix_fmt_)
Definition: filters.h:252
CUDAColorspaceContext::hwctx
AVCUDADeviceContext * hwctx
Definition: vf_colorspace_cuda.c:53
src
#define src
Definition: vp8dsp.c:248
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:3021
options
static const AVOption options[]
Definition: vf_colorspace_cuda.c:389