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 "internal.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  CUDAColorspaceContext* s = ctx->priv;
155  AVHWFramesContext* in_frames_ctx;
156 
157  int ret;
158 
159  if (!ctx->inputs[0]->hw_frames_ctx) {
160  av_log(ctx, AV_LOG_ERROR, "No hw context provided on input\n");
161  return AVERROR(EINVAL);
162  }
163 
164  in_frames_ctx = (AVHWFramesContext*)ctx->inputs[0]->hw_frames_ctx->data;
165  s->pix_fmt = in_frames_ctx->sw_format;
166 
167  if (!format_is_supported(s->pix_fmt)) {
168  av_log(ctx, AV_LOG_ERROR, "Unsupported pixel format: %s\n",
169  av_get_pix_fmt_name(s->pix_fmt));
170  return AVERROR(EINVAL);
171  }
172 
173  if ((AVCOL_RANGE_MPEG != s->range) && (AVCOL_RANGE_JPEG != s->range)) {
174  av_log(ctx, AV_LOG_ERROR, "Unsupported color range\n");
175  return AVERROR(EINVAL);
176  }
177 
178  s->num_planes = av_pix_fmt_count_planes(s->pix_fmt);
179 
180  ret = init_hwframe_ctx(s, in_frames_ctx->device_ref, width, height);
181  if (ret < 0)
182  return ret;
183 
184  ctx->outputs[0]->hw_frames_ctx = av_buffer_ref(s->frames_ctx);
185  if (!ctx->outputs[0]->hw_frames_ctx)
186  return AVERROR(ENOMEM);
187 
188  return 0;
189 }
190 
192 {
193  CUDAColorspaceContext* s = ctx->priv;
194  CUcontext dummy, cuda_ctx = s->hwctx->cuda_ctx;
195  CudaFunctions* cu = s->hwctx->internal->cuda_dl;
196  int ret;
197 
198  extern const unsigned char ff_vf_colorspace_cuda_ptx_data[];
199  extern const unsigned int ff_vf_colorspace_cuda_ptx_len;
200 
201  ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx));
202  if (ret < 0)
203  return ret;
204 
205  ret = ff_cuda_load_module(ctx, s->hwctx, &s->cu_module,
206  ff_vf_colorspace_cuda_ptx_data,
207  ff_vf_colorspace_cuda_ptx_len);
208  if (ret < 0)
209  goto fail;
210 
211  ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_convert[AVCOL_RANGE_MPEG], s->cu_module, "to_mpeg_cuda"));
212  if (ret < 0)
213  goto fail;
214 
215  ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_convert[AVCOL_RANGE_JPEG], s->cu_module, "to_jpeg_cuda"));
216  if (ret < 0)
217  goto fail;
218 
219 fail:
220  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
221  return ret;
222 }
223 
225 {
226  AVFilterContext* ctx = outlink->src;
227  AVFilterLink* inlink = outlink->src->inputs[0];
228  CUDAColorspaceContext* s = ctx->priv;
229  AVHWFramesContext* frames_ctx =
230  (AVHWFramesContext*)inlink->hw_frames_ctx->data;
231  AVCUDADeviceContext* device_hwctx = frames_ctx->device_ctx->hwctx;
232  int ret;
233 
234  s->hwctx = device_hwctx;
235  s->cu_stream = s->hwctx->stream;
236 
237  outlink->w = inlink->w;
238  outlink->h = inlink->h;
239 
241  if (ret < 0)
242  return ret;
243 
244  if (inlink->sample_aspect_ratio.num) {
245  outlink->sample_aspect_ratio = av_mul_q(
246  (AVRational){outlink->h * inlink->w, outlink->w * inlink->h},
247  inlink->sample_aspect_ratio);
248  } else {
249  outlink->sample_aspect_ratio = inlink->sample_aspect_ratio;
250  }
251 
253  if (ret < 0)
254  return ret;
255 
256  return ret;
257 }
258 
260 {
261  CUDAColorspaceContext* s = ctx->priv;
262  CudaFunctions* cu = s->hwctx->internal->cuda_dl;
263  CUcontext dummy, cuda_ctx = s->hwctx->cuda_ctx;
264  int ret;
265 
266  ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx));
267  if (ret < 0)
268  return ret;
269 
270  out->color_range = s->range;
271 
272  for (int i = 0; i < s->num_planes; i++) {
273  int width = in->width, height = in->height, comp_id = (i > 0);
274 
275  switch (s->pix_fmt) {
276  case AV_PIX_FMT_YUV444P:
277  break;
278  case AV_PIX_FMT_YUV420P:
279  width = comp_id ? in->width / 2 : in->width;
280  /* fall-through */
281  case AV_PIX_FMT_NV12:
282  height = comp_id ? in->height / 2 : in->height;
283  break;
284  default:
285  av_log(ctx, AV_LOG_ERROR, "Unsupported pixel format: %s\n",
286  av_get_pix_fmt_name(s->pix_fmt));
287  return AVERROR(EINVAL);
288  }
289 
290  if (!s->cu_convert[out->color_range]) {
291  av_log(ctx, AV_LOG_ERROR, "Unsupported color range\n");
292  return AVERROR(EINVAL);
293  }
294 
295  if (in->color_range != out->color_range) {
296  void* args[] = {&in->data[i], &out->data[i], &in->linesize[i],
297  &comp_id};
298  ret = CHECK_CU(cu->cuLaunchKernel(
299  s->cu_convert[out->color_range], DIV_UP(width, BLOCKX),
300  DIV_UP(height, BLOCKY), 1, BLOCKX, BLOCKY, 1, 0, s->cu_stream,
301  args, NULL));
302  } else {
303  ret = av_hwframe_transfer_data(out, in, 0);
304  if (ret < 0)
305  return ret;
306  }
307  }
308 
309  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
310  return ret;
311 }
312 
314 {
315  CUDAColorspaceContext* s = ctx->priv;
316  AVFilterLink* outlink = ctx->outputs[0];
317  AVFrame* src = in;
318  int ret;
319 
320  ret = conv_cuda_convert(ctx, s->own_frame, src);
321  if (ret < 0)
322  return ret;
323 
324  src = s->own_frame;
325  ret = av_hwframe_get_buffer(src->hw_frames_ctx, s->tmp_frame, 0);
326  if (ret < 0)
327  return ret;
328 
329  av_frame_move_ref(out, s->own_frame);
330  av_frame_move_ref(s->own_frame, s->tmp_frame);
331 
332  s->own_frame->width = outlink->w;
333  s->own_frame->height = outlink->h;
334 
335  ret = av_frame_copy_props(out, in);
336  if (ret < 0)
337  return ret;
338 
339  return 0;
340 }
341 
343 {
344  AVFilterContext* ctx = link->dst;
345  CUDAColorspaceContext* s = ctx->priv;
346  AVFilterLink* outlink = ctx->outputs[0];
347  CudaFunctions* cu = s->hwctx->internal->cuda_dl;
348 
349  AVFrame* out = NULL;
350  CUcontext dummy;
351  int ret = 0;
352 
353  out = av_frame_alloc();
354  if (!out) {
355  ret = AVERROR(ENOMEM);
356  goto fail;
357  }
358 
359  ret = CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx));
360  if (ret < 0)
361  goto fail;
362 
363  ret = cudacolorspace_conv(ctx, out, in);
364 
365  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
366  if (ret < 0)
367  goto fail;
368 
369  av_reduce(&out->sample_aspect_ratio.num, &out->sample_aspect_ratio.den,
370  (int64_t)in->sample_aspect_ratio.num * outlink->h * link->w,
371  (int64_t)in->sample_aspect_ratio.den * outlink->w * link->h,
372  INT_MAX);
373 
374  av_frame_free(&in);
375  return ff_filter_frame(outlink, out);
376 fail:
377  av_frame_free(&in);
378  av_frame_free(&out);
379  return ret;
380 }
381 
382 #define OFFSET(x) offsetof(CUDAColorspaceContext, x)
383 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
384 static const AVOption options[] = {
385  {"range", "Output video range", OFFSET(range), AV_OPT_TYPE_INT, { .i64 = AVCOL_RANGE_UNSPECIFIED }, AVCOL_RANGE_UNSPECIFIED, AVCOL_RANGE_NB - 1, FLAGS, .unit = "range"},
386  {"tv", "Limited range", 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_RANGE_MPEG }, 0, 0, FLAGS, .unit = "range"},
387  {"mpeg", "Limited range", 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_RANGE_MPEG }, 0, 0, FLAGS, .unit = "range"},
388  {"pc", "Full range", 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_RANGE_JPEG }, 0, 0, FLAGS, .unit = "range"},
389  {"jpeg", "Full range", 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_RANGE_JPEG }, 0, 0, FLAGS, .unit = "range"},
390  {NULL},
391 };
392 
394  .class_name = "colorspace_cuda",
395  .item_name = av_default_item_name,
396  .option = options,
397  .version = LIBAVUTIL_VERSION_INT,
398 };
399 
401  {
402  .name = "default",
403  .type = AVMEDIA_TYPE_VIDEO,
404  .filter_frame = cudacolorspace_filter_frame,
405  },
406 };
407 
409  {
410  .name = "default",
411  .type = AVMEDIA_TYPE_VIDEO,
412  .config_props = cudacolorspace_config_props,
413  },
414 };
415 
417  .name = "colorspace_cuda",
418  .description = NULL_IF_CONFIG_SMALL("CUDA accelerated video color converter"),
419 
420  .init = cudacolorspace_init,
421  .uninit = cudacolorspace_uninit,
422 
423  .priv_size = sizeof(CUDAColorspaceContext),
424  .priv_class = &cudacolorspace_class,
425 
428 
430 
431  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
432 };
AVHWDeviceContext::hwctx
void * hwctx
The format-specific data, allocated and freed by libavutil along with this context.
Definition: hwcontext.h:93
AVFrame::color_range
enum AVColorRange color_range
MPEG vs JPEG YUV range.
Definition: frame.h:656
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: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:343
ff_filter_frame
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition: avfilter.c:1007
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:210
cudacolorspace_config_props
static av_cold int cudacolorspace_config_props(AVFilterLink *outlink)
Definition: vf_colorspace_cuda.c:224
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: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:100
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:337
AVFrame
This structure describes decoded (raw) audio or video data.
Definition: frame.h:340
pixdesc.h
CUDAColorspaceContext::num_planes
int num_planes
Definition: vf_colorspace_cuda.c:66
AVFrame::width
int width
Definition: frame.h:412
AVCOL_RANGE_JPEG
@ AVCOL_RANGE_JPEG
Full range content.
Definition: pixfmt.h:687
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:252
AVOption
AVOption.
Definition: opt.h:251
OFFSET
#define OFFSET(x)
Definition: vf_colorspace_cuda.c:382
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:170
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:230
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:361
FLAGS
#define FLAGS
Definition: vf_colorspace_cuda.c:383
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:3008
format_is_supported
static int format_is_supported(enum AVPixelFormat fmt)
Definition: vf_colorspace_cuda.c:142
fail
#define fail()
Definition: checkasm.h:179
dummy
int dummy
Definition: motion.c:66
AVCOL_RANGE_NB
@ AVCOL_RANGE_NB
Not part of ABI.
Definition: pixfmt.h:688
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:33
av_frame_alloc
AVFrame * av_frame_alloc(void)
Allocate an AVFrame and set its fields to default values.
Definition: frame.c:88
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: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:230
width
#define width
s
#define s(width, name)
Definition: cbs_vp9.c:198
CUDAColorspaceContext::own_frame
AVFrame * own_frame
Definition: vf_colorspace_cuda.c:55
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:73
FILTER_INPUTS
#define FILTER_INPUTS(array)
Definition: internal.h:182
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:223
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:736
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:142
AVFilterContext::inputs
AVFilterLink ** inputs
array of pointers to input links
Definition: avfilter.h:417
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:313
AVCOL_RANGE_UNSPECIFIED
@ AVCOL_RANGE_UNSPECIFIED
Definition: pixfmt.h:653
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:106
supported_formats
static enum AVPixelFormat supported_formats[]
Definition: vf_colorspace_cuda.c:38
range
enum AVColorRange range
Definition: mediacodec_wrapper.c:2646
BLOCKX
#define BLOCKX
Definition: vf_colorspace_cuda.c:45
height
#define height
internal.h
FILTER_SINGLE_PIXFMT
#define FILTER_SINGLE_PIXFMT(pix_fmt_)
Definition: internal.h:172
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:255
cudacolorspace_inputs
static const AVFilterPad cudacolorspace_inputs[]
Definition: vf_colorspace_cuda.c:400
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:649
av_frame_unref
void av_frame_unref(AVFrame *frame)
Unreference all the buffers referenced by frame and reset the frame fields.
Definition: frame.c:622
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: internal.h:39
AVCOL_RANGE_MPEG
@ AVCOL_RANGE_MPEG
Narrow or limited range content.
Definition: pixfmt.h:670
ff_vf_colorspace_cuda
const AVFilter ff_vf_colorspace_cuda
Definition: vf_colorspace_cuda.c:416
AVFilter
Filter definition.
Definition: avfilter.h:166
AVHWFramesContext
This struct describes a set or pool of "hardware" frames (i.e.
Definition: hwcontext.h:125
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:71
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:150
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:447
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:451
AVFrame::height
int height
Definition: frame.h:412
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:342
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: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:409
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:259
FFALIGN
#define FFALIGN(x, a)
Definition: macros.h:78
FILTER_OUTPUTS
#define FILTER_OUTPUTS(array)
Definition: internal.h:183
src
INIT_CLIP pixel * src
Definition: h264pred_template.c:418
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:385
av_log
#define av_log(a,...)
Definition: tableprint_vlc.h:27
cudacolorspace_class
static const AVClass cudacolorspace_class
Definition: vf_colorspace_cuda.c:393
cudacolorspace_outputs
static const AVFilterPad cudacolorspace_outputs[]
Definition: vf_colorspace_cuda.c:408
cudacolorspace_load_functions
static av_cold int cudacolorspace_load_functions(AVFilterContext *ctx)
Definition: vf_colorspace_cuda.c:191
AVColorRange
AVColorRange
Visual content value range.
Definition: pixfmt.h:652
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:510
AV_OPT_TYPE_CONST
@ AV_OPT_TYPE_CONST
Definition: opt.h:234
CUDAColorspaceContext::hwctx
AVCUDADeviceContext * hwctx
Definition: vf_colorspace_cuda.c:53
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:2888
options
static const AVOption options[]
Definition: vf_colorspace_cuda.c:384