FFmpeg: libavfilter/vf_chromakey_cuda.c Source File

FFmpeg

[フレーム]

libavfilter

vf_chromakey_cuda.c

Go to the documentation of this file.

1 /*

3 *

4 * This file is part of FFmpeg.

5 *

6 * FFmpeg is free software; you can redistribute it and/or

7 * modify it under the terms of the GNU Lesser General Public

8 * License as published by the Free Software Foundation; either

9 * version 2.1 of the License, or (at your option) any later version.

10 *

11 * FFmpeg is distributed in the hope that it will be useful,

12 * but WITHOUT ANY WARRANTY; without even the implied warranty of

13 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU

14 * Lesser General Public License for more details.

15 *

16 * You should have received a copy of the GNU Lesser General Public

17 * License along with FFmpeg; if not, write to the Free Software

18 * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA

19 */

21 #include <float.h>

22 #include <stdio.h>

24 #include "libavutil/common.h"

25 #include "libavutil/hwcontext.h"

26 #include "libavutil/hwcontext_cuda_internal.h"

27 #include "libavutil/cuda_check.h"

28 #include "libavutil/internal.h"

29 #include "libavutil/opt.h"

30 #include "libavutil/pixdesc.h"

32 #include "avfilter.h"

33 #include "filters.h"

34 #include "cuda/load_helper.h"

36 static const enum AVPixelFormat supported_formats[] = {

37 AV_PIX_FMT_YUV420P,

38 AV_PIX_FMT_NV12,

39 AV_PIX_FMT_YUVA420P

40 };

42 #define DIV_UP(a, b) (((a) + (b)-1) / (b))

43 #define BLOCKX 32

44 #define BLOCKY 16

45 #define CHECK_CU(x) FF_CUDA_CHECK_DL(ctx, s->hwctx->internal->cuda_dl, x)

47 typedef struct ChromakeyCUDAContext {

48 const AVClass *class;

50 AVCUDADeviceContext *hwctx;

52 enum AVPixelFormat in_fmt, out_fmt;

53 const AVPixFmtDescriptor *in_desc, *out_desc;

54 int in_planes, out_planes;

55 int in_plane_depths[4];

56 int in_plane_channels[4];

58 uint8_t chromakey_rgba[4];

59 uint16_t chromakey_uv[2];

60 int is_yuv;

61 float similarity;

62 float blend;

64 AVBufferRef *frames_ctx;

65 AVFrame *frame;

66 AVFrame *tmp_frame;

68 CUcontext cu_ctx;

69 CUmodule cu_module;

70 CUfunction cu_func;

71 CUfunction cu_func_uv;

72 CUstream cu_stream;

73 } ChromakeyCUDAContext;

75 static av_cold int cudachromakey_init(AVFilterContext *ctx)

76 {

77 ChromakeyCUDAContext *s = ctx->priv;

79 s->frame = av_frame_alloc();

80 if (!s->frame)

81 return AVERROR(ENOMEM);

83 s->tmp_frame = av_frame_alloc();

84 if (!s->tmp_frame)

85 return AVERROR(ENOMEM);

87 return 0;

88 }

90 static av_cold void cudachromakey_uninit(AVFilterContext *ctx)

91 {

92 ChromakeyCUDAContext *s = ctx->priv;

94 if (s->hwctx && s->cu_module)

95 {

96 CudaFunctions *cu = s->hwctx->internal->cuda_dl;

97 CUcontext context;

99 CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx));

100 CHECK_CU(cu->cuModuleUnload(s->cu_module));

101 s->cu_module = NULL;

102 CHECK_CU(cu->cuCtxPopCurrent(&context));

103 }

104

105 av_frame_free(&s->frame);

106 av_buffer_unref(&s->frames_ctx);

107 av_frame_free(&s->tmp_frame);

108 }

109

110 static av_cold int init_hwframe_ctx(ChromakeyCUDAContext *s, AVBufferRef *device_ctx, int width, int height)

111 {

112 AVBufferRef *out_ref = NULL;

113 AVHWFramesContext *out_ctx;

114 int ret;

115

116 out_ref = av_hwframe_ctx_alloc(device_ctx);

117 if (!out_ref)

118 return AVERROR(ENOMEM);

119 out_ctx = (AVHWFramesContext *)out_ref->data;

120

121 out_ctx->format = AV_PIX_FMT_CUDA;

122 out_ctx->sw_format = s->out_fmt;

123 out_ctx->width = width;

124 out_ctx->height = height;

125

126 ret = av_hwframe_ctx_init(out_ref);

127 if (ret < 0)

128 goto fail;

129

130 av_frame_unref(s->frame);

131 ret = av_hwframe_get_buffer(out_ref, s->frame, 0);

132 if (ret < 0)

133 goto fail;

134

135 av_buffer_unref(&s->frames_ctx);

136 s->frames_ctx = out_ref;

137

138 return 0;

139 fail:

140 av_buffer_unref(&out_ref);

141 return ret;

142 }

143

144 static int format_is_supported(enum AVPixelFormat fmt)

145 {

146 int i;

147

148 for (i = 0; i < FF_ARRAY_ELEMS(supported_formats); i++)

149 if (supported_formats[i] == fmt)

150 return 1;

151 return 0;

152 }

153

154 static av_cold void set_format_info(AVFilterContext *ctx, enum AVPixelFormat in_format, enum AVPixelFormat out_format)

155 {

156 ChromakeyCUDAContext *s = ctx->priv;

157 int i, p, d;

158

159 s->in_fmt = in_format;

160 s->out_fmt = out_format;

161

162 s->in_desc = av_pix_fmt_desc_get(s->in_fmt);

163 s->out_desc = av_pix_fmt_desc_get(s->out_fmt);

164 s->in_planes = av_pix_fmt_count_planes(s->in_fmt);

165 s->out_planes = av_pix_fmt_count_planes(s->out_fmt);

166

167 // find maximum step of each component of each plane

168 // For our subset of formats, this should accurately tell us how many channels CUDA needs

169 // i.e. 1 for Y plane, 2 for UV plane of NV12, 4 for single plane of RGB0 formats

170

171 for (i = 0; i < s->in_desc->nb_components; i++)

172 {

173 d = (s->in_desc->comp[i].depth + 7) / 8;

174 p = s->in_desc->comp[i].plane;

175 s->in_plane_channels[p] = FFMAX(s->in_plane_channels[p], s->in_desc->comp[i].step / d);

176

177 s->in_plane_depths[p] = s->in_desc->comp[i].depth;

178 }

179 }

180

181 static av_cold int init_processing_chain(AVFilterContext *ctx, int width, int height)

182 {

183 FilterLink *inl = ff_filter_link(ctx->inputs[0]);

184 FilterLink *outl = ff_filter_link(ctx->outputs[0]);

185 ChromakeyCUDAContext *s = ctx->priv;

186 AVHWFramesContext *in_frames_ctx;

187 int ret;

188

189 /* check that we have a hw context */

190 if (!inl->hw_frames_ctx) {

191 av_log(ctx, AV_LOG_ERROR, "No hw context provided on input\n");

192 return AVERROR(EINVAL);

193 }

194 in_frames_ctx = (AVHWFramesContext *)inl->hw_frames_ctx->data;

195

196 if (!format_is_supported(in_frames_ctx->sw_format))

197 {

198 av_log(ctx, AV_LOG_ERROR, "Unsupported format: %s\n", av_get_pix_fmt_name(in_frames_ctx->sw_format));

199 return AVERROR(ENOSYS);

200 }

201

202 set_format_info(ctx, in_frames_ctx->sw_format, AV_PIX_FMT_YUVA420P);

203

204 ret = init_hwframe_ctx(s, in_frames_ctx->device_ref, width, height);

205 if (ret < 0)

206 return ret;

207

208 outl->hw_frames_ctx = av_buffer_ref(s->frames_ctx);

209 if (!outl->hw_frames_ctx)

210 return AVERROR(ENOMEM);

211

212 return 0;

213 }

214

215 static av_cold int cudachromakey_load_functions(AVFilterContext *ctx)

216 {

217 ChromakeyCUDAContext *s = ctx->priv;

218 CUcontext context, cuda_ctx = s->hwctx->cuda_ctx;

219 CudaFunctions *cu = s->hwctx->internal->cuda_dl;

220 int ret;

221

222 extern const unsigned char ff_vf_chromakey_cuda_ptx_data[];

223 extern const unsigned int ff_vf_chromakey_cuda_ptx_len;

224

225 ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx));

226 if (ret < 0)

227 return ret;

228

229 ret = ff_cuda_load_module(ctx, s->hwctx, &s->cu_module,

230 ff_vf_chromakey_cuda_ptx_data, ff_vf_chromakey_cuda_ptx_len);

231 if (ret < 0)

232 goto fail;

233

234 ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func, s->cu_module, "Process_uchar"));

235 if (ret < 0)

236 {

237 av_log(ctx, AV_LOG_FATAL, "Failed loading Process_uchar\n");

238 goto fail;

239 }

240

241 ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uv, s->cu_module, "Process_uchar2"));

242 if (ret < 0)

243 {

244 av_log(ctx, AV_LOG_FATAL, "Failed loading Process_uchar2\n");

245 goto fail;

246 }

247

248 fail:

249 CHECK_CU(cu->cuCtxPopCurrent(&context));

250

251 return ret;

252 }

253

254 #define FIXNUM(x) lrint((x) * (1 << 10))

255 #define RGB_TO_U(rgb) (((-FIXNUM(0.16874) * rgb[0] - FIXNUM(0.33126) * rgb[1] + FIXNUM(0.50000) * rgb[2] + (1 << 9) - 1) >> 10) + 128)

256 #define RGB_TO_V(rgb) (((FIXNUM(0.50000) * rgb[0] - FIXNUM(0.41869) * rgb[1] - FIXNUM(0.08131) * rgb[2] + (1 << 9) - 1) >> 10) + 128)

257

258 static av_cold int cudachromakey_config_props(AVFilterLink *outlink)

259 {

260 AVFilterContext *ctx = outlink->src;

261 AVFilterLink *inlink = outlink->src->inputs[0];

262 FilterLink *inl = ff_filter_link(inlink);

263 ChromakeyCUDAContext *s = ctx->priv;

264 AVHWFramesContext *frames_ctx;

265 AVCUDADeviceContext *device_hwctx;

266 int ret;

267

268 if (s->is_yuv) {

269 s->chromakey_uv[0] = s->chromakey_rgba[1];

270 s->chromakey_uv[1] = s->chromakey_rgba[2];

271 } else {

272 s->chromakey_uv[0] = RGB_TO_U(s->chromakey_rgba);

273 s->chromakey_uv[1] = RGB_TO_V(s->chromakey_rgba);

274 }

275

276 ret = init_processing_chain(ctx, inlink->w, inlink->h);

277 if (ret < 0)

278 return ret;

279

280 frames_ctx = (AVHWFramesContext *)inl->hw_frames_ctx->data;

281 device_hwctx = frames_ctx->device_ctx->hwctx;

282

283 s->hwctx = device_hwctx;

284 s->cu_stream = s->hwctx->stream;

285

286 outlink->sample_aspect_ratio = inlink->sample_aspect_ratio;

287

288 ret = cudachromakey_load_functions(ctx);

289 if (ret < 0)

290 return ret;

291

292 return 0;

293 }

294

295 static int call_cuda_kernel(AVFilterContext *ctx, CUfunction func,

296 CUtexObject src_tex[3], AVFrame *out_frame,

297 int width, int height, int pitch,

298 int width_uv, int height_uv, int pitch_uv,

299 float u_key, float v_key, float similarity,

300 float blend)

301 {

302 ChromakeyCUDAContext *s = ctx->priv;

303 CudaFunctions *cu = s->hwctx->internal->cuda_dl;

304

305 CUdeviceptr dst_devptr[4] = {

306 (CUdeviceptr)out_frame->data[0], (CUdeviceptr)out_frame->data[1],

307 (CUdeviceptr)out_frame->data[2], (CUdeviceptr)out_frame->data[3]

308 };

309

310 void *args_uchar[] = {

311 &src_tex[0], &src_tex[1], &src_tex[2],

312 &dst_devptr[0], &dst_devptr[1], &dst_devptr[2], &dst_devptr[3],

313 &width, &height, &pitch,

314 &width_uv, &height_uv, &pitch_uv,

315 &u_key, &v_key, &similarity, &blend

316 };

317

318 return CHECK_CU(cu->cuLaunchKernel(func,

319 DIV_UP(width, BLOCKX), DIV_UP(height, BLOCKY), 1,

320 BLOCKX, BLOCKY, 1, 0, s->cu_stream, args_uchar, NULL));

321 }

322

323 static int cudachromakey_process_internal(AVFilterContext *ctx,

324 AVFrame *out, AVFrame *in)

325 {

326 ChromakeyCUDAContext *s = ctx->priv;

327 CudaFunctions *cu = s->hwctx->internal->cuda_dl;

328 CUcontext context, cuda_ctx = s->hwctx->cuda_ctx;

329 int i, ret;

330

331 CUtexObject tex[3] = {0, 0, 0};

332

333 ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx));

334 if (ret < 0)

335 return ret;

336

337 for (i = 0; i < s->in_planes; i++)

338 {

339 CUDA_TEXTURE_DESC tex_desc = {

340 .filterMode = CU_TR_FILTER_MODE_LINEAR,

341 .flags = 0, // CU_TRSF_READ_AS_INTEGER to get raw ints instead of normalized floats from tex2D

342 };

343

344 CUDA_RESOURCE_DESC res_desc = {

345 .resType = CU_RESOURCE_TYPE_PITCH2D,

346 .res.pitch2D.format = CU_AD_FORMAT_UNSIGNED_INT8,

347 .res.pitch2D.numChannels = s->in_plane_channels[i],

348 .res.pitch2D.pitchInBytes = in->linesize[i],

349 .res.pitch2D.devPtr = (CUdeviceptr)in->data[i],

350 };

351

352 if (i == 1 || i == 2)

353 {

354 res_desc.res.pitch2D.width = AV_CEIL_RSHIFT(in->width, s->in_desc->log2_chroma_w);

355 res_desc.res.pitch2D.height = AV_CEIL_RSHIFT(in->height, s->in_desc->log2_chroma_h);

356 }

357 else

358 {

359 res_desc.res.pitch2D.width = in->width;

360 res_desc.res.pitch2D.height = in->height;

361 }

362

363 ret = CHECK_CU(cu->cuTexObjectCreate(&tex[i], &res_desc, &tex_desc, NULL));

364 if (ret < 0)

365 goto exit;

366 }

367

368 ret = call_cuda_kernel(ctx, (s->in_plane_channels[1] > 1) ? s->cu_func_uv : s->cu_func,

369 tex, out,

370 out->width, out->height, out->linesize[0],

371 AV_CEIL_RSHIFT(out->width, s->out_desc->log2_chroma_w),

372 AV_CEIL_RSHIFT(out->height, s->out_desc->log2_chroma_h),

373 out->linesize[1],

374 s->chromakey_uv[0], s->chromakey_uv[1],

375 s->similarity, s->blend);

376 if (ret < 0)

377 goto exit;

378

379 exit:

380 for (i = 0; i < s->in_planes; i++)

381 if (tex[i])

382 CHECK_CU(cu->cuTexObjectDestroy(tex[i]));

383

384 CHECK_CU(cu->cuCtxPopCurrent(&context));

385

386 return ret;

387 }

388

389 static int cudachromakey_process(AVFilterContext *ctx, AVFrame *out, AVFrame *in)

390 {

391 ChromakeyCUDAContext *s = ctx->priv;

392 AVFrame *src = in;

393 int ret;

394

395 ret = cudachromakey_process_internal(ctx, s->frame, src);

396 if (ret < 0)

397 return ret;

398

399 src = s->frame;

400 ret = av_hwframe_get_buffer(src->hw_frames_ctx, s->tmp_frame, 0);

401 if (ret < 0)

402 return ret;

403

404 av_frame_move_ref(out, s->frame);

405 av_frame_move_ref(s->frame, s->tmp_frame);

406

407 ret = av_frame_copy_props(out, in);

408 if (ret < 0)

409 return ret;

410

411 return 0;

412 }

413

414 static int cudachromakey_filter_frame(AVFilterLink *link, AVFrame *in)

415 {

416 AVFilterContext *ctx = link->dst;

417 ChromakeyCUDAContext *s = ctx->priv;

418 AVFilterLink *outlink = ctx->outputs[0];

419 CudaFunctions *cu = s->hwctx->internal->cuda_dl;

420

421 AVFrame *out = NULL;

422 CUcontext context;

423 int ret = 0;

424

425 out = av_frame_alloc();

426 if (!out)

427 {

428 ret = AVERROR(ENOMEM);

429 goto fail;

430 }

431

432 ret = CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx));

433 if (ret < 0)

434 goto fail;

435

436 ret = cudachromakey_process(ctx, out, in);

437

438 CHECK_CU(cu->cuCtxPopCurrent(&context));

439 if (ret < 0)

440 goto fail;

441

442 av_frame_free(&in);

443 return ff_filter_frame(outlink, out);

444 fail:

445 av_frame_free(&in);

446 av_frame_free(&out);

447 return ret;

448 }

449

450 #define OFFSET(x) offsetof(ChromakeyCUDAContext, x)

451 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)

452 static const AVOption options[] = {

453 {"color", "set the chromakey key color", OFFSET(chromakey_rgba), AV_OPT_TYPE_COLOR, {.str = "black"}, 0, 0, FLAGS},

454 {"similarity", "set the chromakey similarity value", OFFSET(similarity), AV_OPT_TYPE_FLOAT, {.dbl = 0.01}, 0.01, 1.0, FLAGS},

455 {"blend", "set the chromakey key blend value", OFFSET(blend), AV_OPT_TYPE_FLOAT, {.dbl = 0.0}, 0.0, 1.0, FLAGS},

456 {"yuv", "color parameter is in yuv instead of rgb", OFFSET(is_yuv), AV_OPT_TYPE_BOOL, {.i64 = 0}, 0, 1, FLAGS},

457 {NULL},

458 };

459

460 static const AVClass cudachromakey_class = {

461 .class_name = "cudachromakey",

462 .item_name = av_default_item_name,

463 .option = options,

464 .version = LIBAVUTIL_VERSION_INT,

465 };

466

467 static const AVFilterPad cudachromakey_inputs[] = {

468 {

469 .name = "default",

470 .type = AVMEDIA_TYPE_VIDEO,

471 .filter_frame = cudachromakey_filter_frame,

472 },

473 };

474

475 static const AVFilterPad cudachromakey_outputs[] = {

476 {

477 .name = "default",

478 .type = AVMEDIA_TYPE_VIDEO,

479 .config_props = cudachromakey_config_props,

480 },

481 };

482

483 const AVFilter ff_vf_chromakey_cuda = {

484 .name = "chromakey_cuda",

485 .description = NULL_IF_CONFIG_SMALL("GPU accelerated chromakey filter"),

486

487 .init = cudachromakey_init,

488 .uninit = cudachromakey_uninit,

489

490 .priv_size = sizeof(ChromakeyCUDAContext),

491 .priv_class = &cudachromakey_class,

492

493 FILTER_INPUTS(cudachromakey_inputs),

494 FILTER_OUTPUTS(cudachromakey_outputs),

495

496 FILTER_SINGLE_PIXFMT(AV_PIX_FMT_CUDA),

497

498 .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,

499 };

func

int(* func)(AVBPrint *dst, const char *in, const char *arg)

Definition: jacosubdec.c:68

ChromakeyCUDAContext

Definition: vf_chromakey_cuda.c:47

AVHWDeviceContext::hwctx

void * hwctx

The format-specific data, allocated and freed by libavutil along with this context.

Definition: hwcontext.h:85

call_cuda_kernel

static int call_cuda_kernel(AVFilterContext *ctx, CUfunction func, CUtexObject src_tex[3], AVFrame *out_frame, int width, int height, int pitch, int width_uv, int height_uv, int pitch_uv, float u_key, float v_key, float similarity, float blend)

Definition: vf_chromakey_cuda.c:295

AV_PIX_FMT_CUDA

@ AV_PIX_FMT_CUDA

HW acceleration through CUDA.

Definition: pixfmt.h:260

AVPixelFormat

Pixel format.

Definition: pixfmt.h:71

supported_formats

static enum AVPixelFormat supported_formats[]

Definition: vf_chromakey_cuda.c:36

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

ChromakeyCUDAContext::in_desc

const AVPixFmtDescriptor * in_desc

Definition: vf_chromakey_cuda.c:53

ff_filter_frame

int ff_filter_frame(AVFilterLink *link, AVFrame *frame)

Send a frame of data to the next filter.

Definition: avfilter.c:1023

av_pix_fmt_desc_get

const AVPixFmtDescriptor * av_pix_fmt_desc_get(enum AVPixelFormat pix_fmt)

Definition: pixdesc.c:2965

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

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

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

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

AVFrame::width

int width

Definition: frame.h:461

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

ChromakeyCUDAContext::cu_func_uv

CUfunction cu_func_uv

Definition: vf_chromakey_cuda.c:71

ChromakeyCUDAContext::chromakey_rgba

uint8_t chromakey_rgba[4]

Definition: vf_chromakey_cuda.c:58

float.h

FFMAX

#define FFMAX(a, b)

Definition: macros.h:47

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

AVHWFramesContext::width

int width

The allocated dimensions of the frames in this pool.

Definition: hwcontext.h:217

AVFilterLink

A link between two filters.

Definition: avfilter.h:557

DIV_UP

#define DIV_UP(a, b)

Definition: vf_chromakey_cuda.c:42

ChromakeyCUDAContext::tmp_frame

AVFrame * tmp_frame

Definition: vf_chromakey_cuda.c:66

AVFrame::data

uint8_t * data[AV_NUM_DATA_POINTERS]

pointer to the picture/channel planes.

Definition: frame.h:410

FilterLink

Link properties exposed to filter code, but not external callers.

Definition: filters.h:127

av_pix_fmt_count_planes

int av_pix_fmt_count_planes(enum AVPixelFormat pix_fmt)

Definition: pixdesc.c:3005

fail

#define fail()

Definition: checkasm.h:188

cudachromakey_load_functions

static av_cold int cudachromakey_load_functions(AVFilterContext *ctx)

Definition: vf_chromakey_cuda.c:215

ChromakeyCUDAContext::is_yuv

int is_yuv

Definition: vf_chromakey_cuda.c:60

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

ChromakeyCUDAContext::cu_ctx

CUcontext cu_ctx

Definition: vf_chromakey_cuda.c:68

FLAGS

#define FLAGS

Definition: vf_chromakey_cuda.c:451

BLOCKY

#define BLOCKY

Definition: vf_chromakey_cuda.c:44

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:217

ChromakeyCUDAContext::frames_ctx

AVBufferRef * frames_ctx

Definition: vf_chromakey_cuda.c:64

#define s(width, name)

Definition: cbs_vp9.c:198

AV_PIX_FMT_YUVA420P

@ AV_PIX_FMT_YUVA420P

planar YUV 4:2:0, 20bpp, (1 Cr & Cb sample per 2x2 Y & A samples)

Definition: pixfmt.h:108

AV_CEIL_RSHIFT

#define AV_CEIL_RSHIFT(a, b)

Definition: common.h:60

ff_vf_chromakey_cuda

const AVFilter ff_vf_chromakey_cuda

Definition: vf_chromakey_cuda.c:483

ChromakeyCUDAContext::cu_stream

CUstream cu_stream

Definition: vf_chromakey_cuda.c:72

AVFilterLink::sample_aspect_ratio

AVRational sample_aspect_ratio

agreed upon sample aspect ratio

Definition: avfilter.h:571

ChromakeyCUDAContext::in_plane_depths

int in_plane_depths[4]

Definition: vf_chromakey_cuda.c:55

RGB_TO_U

#define RGB_TO_U(rgb)

Definition: vf_chromakey_cuda.c:255

ChromakeyCUDAContext::blend

float blend

Definition: vf_chromakey_cuda.c:62

filters.h

options

static const AVOption options[]

Definition: vf_chromakey_cuda.c:452

ctx

AVFormatContext * ctx

Definition: movenc.c:49

cudachromakey_uninit

static av_cold void cudachromakey_uninit(AVFilterContext *ctx)

Definition: vf_chromakey_cuda.c:90

OFFSET

#define OFFSET(x)

Definition: vf_chromakey_cuda.c:450

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

Definition: filter_design.txt:23

if(ret)

Definition: filter_design.txt:179

context

it s the only field you need to keep assuming you have a context There is some magic you don t need to care about around this just let it vf default minimum maximum flags name is the option keep it simple and lowercase description are in without and describe what they for example set the foo of the bar offset is the offset of the field in your context

Definition: writing_filters.txt:91

LIBAVUTIL_VERSION_INT

#define LIBAVUTIL_VERSION_INT

Definition: version.h:85

AVClass

Describe the class of an AVClass context structure.

Definition: log.h:66

init_processing_chain

static av_cold int init_processing_chain(AVFilterContext *ctx, int width, int height)

Definition: vf_chromakey_cuda.c:181

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

AVHWFramesContext::device_ref

AVBufferRef * device_ref

A reference to the parent AVHWDeviceContext.

Definition: hwcontext.h:126

ChromakeyCUDAContext::cu_func

CUfunction cu_func

Definition: vf_chromakey_cuda.c:70

AV_OPT_TYPE_COLOR

@ AV_OPT_TYPE_COLOR

Underlying C type is uint8_t[4].

Definition: opt.h:323

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

ChromakeyCUDAContext::hwctx

AVCUDADeviceContext * hwctx

Definition: vf_chromakey_cuda.c:50

ChromakeyCUDAContext::out_fmt

enum AVPixelFormat in_fmt out_fmt

Definition: vf_chromakey_cuda.c:52

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

ChromakeyCUDAContext::out_desc

const AVPixFmtDescriptor * out_desc

Definition: vf_chromakey_cuda.c:53

init_hwframe_ctx

static av_cold int init_hwframe_ctx(ChromakeyCUDAContext *s, AVBufferRef *device_ctx, int width, int height)

Definition: vf_chromakey_cuda.c:110

cudachromakey_init

static av_cold int cudachromakey_init(AVFilterContext *ctx)

Definition: vf_chromakey_cuda.c:75

AVFilterLink::src

AVFilterContext * src

source filter

Definition: avfilter.h:558

ChromakeyCUDAContext::similarity

float similarity

Definition: vf_chromakey_cuda.c:61

ChromakeyCUDAContext::in_plane_channels

int in_plane_channels[4]

Definition: vf_chromakey_cuda.c:56

ChromakeyCUDAContext::out_planes

int out_planes

Definition: vf_chromakey_cuda.c:54

RGB_TO_V

#define RGB_TO_V(rgb)

Definition: vf_chromakey_cuda.c:256

ChromakeyCUDAContext::chromakey_uv

uint16_t chromakey_uv[2]

Definition: vf_chromakey_cuda.c:59

FilterLink::hw_frames_ctx

AVBufferRef * hw_frames_ctx

For hwaccel pixel formats, this should be a reference to the AVHWFramesContext describing the frames.

Definition: filters.h:194

AV_OPT_TYPE_FLOAT

@ AV_OPT_TYPE_FLOAT

Underlying C type is float.

Definition: opt.h:271

ChromakeyCUDAContext::cu_module

CUmodule cu_module

Definition: vf_chromakey_cuda.c:69

#define i(width, name, range_min, range_max)

Definition: cbs_h2645.c:256

internal.h

cudachromakey_class

static const AVClass cudachromakey_class

Definition: vf_chromakey_cuda.c:460

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

set_format_info

static av_cold void set_format_info(AVFilterContext *ctx, enum AVPixelFormat in_format, enum AVPixelFormat out_format)

Definition: vf_chromakey_cuda.c:154

AVFilterPad::name

const char * name

Pad name.

Definition: filters.h:44

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

format_is_supported

static int format_is_supported(enum AVPixelFormat fmt)

Definition: vf_chromakey_cuda.c:144

ret

Definition: filter_design.txt:187

AV_LOG_FATAL

#define AV_LOG_FATAL

Something went wrong and recovery is not possible.

Definition: log.h:174

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

AVHWFramesContext::device_ctx

AVHWDeviceContext * device_ctx

The parent AVHWDeviceContext.

Definition: hwcontext.h:134

cuda_check.h

cudachromakey_outputs

static const AVFilterPad cudachromakey_outputs[]

Definition: vf_chromakey_cuda.c:475

AVFrame::height

int height

Definition: frame.h:461

avfilter.h

cudachromakey_filter_frame

static int cudachromakey_filter_frame(AVFilterLink *link, AVFrame *in)

Definition: vf_chromakey_cuda.c:414

cudachromakey_inputs

static const AVFilterPad cudachromakey_inputs[]

Definition: vf_chromakey_cuda.c:467

BLOCKX

#define BLOCKX

Definition: vf_chromakey_cuda.c:43

AVFilterContext

An instance of a filter.

Definition: avfilter.h:457

ChromakeyCUDAContext::in_planes

int in_planes

Definition: vf_chromakey_cuda.c:54

AVMEDIA_TYPE_VIDEO

@ AVMEDIA_TYPE_VIDEO

Definition: avutil.h:201

AVBufferRef

A reference to a data buffer.

Definition: buffer.h:82

AVPixFmtDescriptor

Descriptor that unambiguously describes how the bits of a pixel are stored in the up to 4 data planes...

Definition: pixdesc.h:69

ChromakeyCUDAContext::frame

AVFrame * frame

Definition: vf_chromakey_cuda.c:65

cudachromakey_process_internal

static int cudachromakey_process_internal(AVFilterContext *ctx, AVFrame *out, AVFrame *in)

Definition: vf_chromakey_cuda.c:323

cudachromakey_config_props

static av_cold int cudachromakey_config_props(AVFilterLink *outlink)

Definition: vf_chromakey_cuda.c:258

CHECK_CU

#define CHECK_CU(x)

Definition: vf_chromakey_cuda.c:45

AV_OPT_TYPE_BOOL

@ AV_OPT_TYPE_BOOL

Underlying C type is int.

Definition: opt.h:327

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

width

#define width

Definition: dsp.h:85

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

FILTER_SINGLE_PIXFMT

#define FILTER_SINGLE_PIXFMT(pix_fmt_)

Definition: filters.h:252

cudachromakey_process

static int cudachromakey_process(AVFilterContext *ctx, AVFrame *out, AVFrame *in)

Definition: vf_chromakey_cuda.c:389

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:2885

Generated on Fri Aug 22 2025 13:59:12 for FFmpeg by doxygen 1.8.17