FFmpeg: libavfilter/vf_thumbnail_cuda.c Source File

FFmpeg

[フレーム]

libavfilter

vf_thumbnail_cuda.c

Go to the documentation of this file.

1 /*

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 */

23 #include "libavutil/hwcontext.h"

24 #include "libavutil/hwcontext_cuda_internal.h"

25 #include "libavutil/cuda_check.h"

26 #include "libavutil/opt.h"

27 #include "libavutil/pixdesc.h"

29 #include "avfilter.h"

30 #include "internal.h"

32 #include "cuda/load_helper.h"

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

36 #define HIST_SIZE (3*256)

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

38 #define BLOCKX 32

39 #define BLOCKY 16

41 static const enum AVPixelFormat supported_formats[] = {

42 AV_PIX_FMT_NV12,

43 AV_PIX_FMT_YUV420P,

44 AV_PIX_FMT_YUV444P,

45 AV_PIX_FMT_P010,

46 AV_PIX_FMT_P016,

47 AV_PIX_FMT_YUV444P16,

48 };

50 struct thumb_frame {

51 AVFrame *buf; ///< cached frame

52 int histogram[HIST_SIZE]; ///< RGB color distribution histogram of the frame

53 };

55 typedef struct ThumbnailCudaContext {

56 const AVClass *class;

57 int n; ///< current frame

58 int n_frames; ///< number of frames for analysis

59 struct thumb_frame *frames; ///< the n_frames frames

60 AVRational tb; ///< copy of the input timebase to ease access

62 AVBufferRef *hw_frames_ctx;

63 AVCUDADeviceContext *hwctx;

65 CUmodule cu_module;

67 CUfunction cu_func_uchar;

68 CUfunction cu_func_uchar2;

69 CUfunction cu_func_ushort;

70 CUfunction cu_func_ushort2;

71 CUstream cu_stream;

73 CUdeviceptr data;

75 } ThumbnailCudaContext;

77 #define OFFSET(x) offsetof(ThumbnailCudaContext, x)

78 #define FLAGS AV_OPT_FLAG_VIDEO_PARAM|AV_OPT_FLAG_FILTERING_PARAM

80 static const AVOption thumbnail_cuda_options[] = {

81 { "n", "set the frames batch size", OFFSET(n_frames), AV_OPT_TYPE_INT, {.i64=100}, 2, INT_MAX, FLAGS },

82 { NULL }

83 };

85 AVFILTER_DEFINE_CLASS(thumbnail_cuda);

87 static av_cold int init(AVFilterContext *ctx)

88 {

89 ThumbnailCudaContext *s = ctx->priv;

91 s->frames = av_calloc(s->n_frames, sizeof(*s->frames));

92 if (!s->frames) {

93 av_log(ctx, AV_LOG_ERROR,

94 "Allocation failure, try to lower the number of frames\n");

95 return AVERROR(ENOMEM);

96 }

97 av_log(ctx, AV_LOG_VERBOSE, "batch size: %d frames\n", s->n_frames);

98 return 0;

99 }

100

101 /**

102 * @brief Compute Sum-square deviation to estimate "closeness".

103 * @param hist color distribution histogram

104 * @param median average color distribution histogram

105 * @return sum of squared errors

106 */

107 static double frame_sum_square_err(const int *hist, const double *median)

108 {

109 int i;

110 double err, sum_sq_err = 0;

111

112 for (i = 0; i < HIST_SIZE; i++) {

113 err = median[i] - (double)hist[i];

114 sum_sq_err += err*err;

115 }

116 return sum_sq_err;

117 }

118

119 static AVFrame *get_best_frame(AVFilterContext *ctx)

120 {

121 AVFrame *picref;

122 ThumbnailCudaContext *s = ctx->priv;

123 int i, j, best_frame_idx = 0;

124 int nb_frames = s->n;

125 double avg_hist[HIST_SIZE] = {0}, sq_err, min_sq_err = -1;

126

127 // average histogram of the N frames

128 for (j = 0; j < FF_ARRAY_ELEMS(avg_hist); j++) {

129 for (i = 0; i < nb_frames; i++)

130 avg_hist[j] += (double)s->frames[i].histogram[j];

131 avg_hist[j] /= nb_frames;

132 }

133

134 // find the frame closer to the average using the sum of squared errors

135 for (i = 0; i < nb_frames; i++) {

136 sq_err = frame_sum_square_err(s->frames[i].histogram, avg_hist);

137 if (i == 0 || sq_err < min_sq_err)

138 best_frame_idx = i, min_sq_err = sq_err;

139 }

140

141 // free and reset everything (except the best frame buffer)

142 for (i = 0; i < nb_frames; i++) {

143 memset(s->frames[i].histogram, 0, sizeof(s->frames[i].histogram));

144 if (i != best_frame_idx)

145 av_frame_free(&s->frames[i].buf);

146 }

147 s->n = 0;

148

149 // raise the chosen one

150 picref = s->frames[best_frame_idx].buf;

151 av_log(ctx, AV_LOG_INFO, "frame id #%d (pts_time=%f) selected "

152 "from a set of %d images\n", best_frame_idx,

153 picref->pts * av_q2d(s->tb), nb_frames);

154 s->frames[best_frame_idx].buf = NULL;

155

156 return picref;

157 }

158

159 static int thumbnail_kernel(AVFilterContext *ctx, CUfunction func, int channels,

160 int *histogram, uint8_t *src_dptr, int src_width, int src_height, int src_pitch, int pixel_size)

161 {

162 int ret;

163 ThumbnailCudaContext *s = ctx->priv;

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

165 CUtexObject tex = 0;

166 void *args[] = { &tex, &histogram, &src_width, &src_height };

167

168 CUDA_TEXTURE_DESC tex_desc = {

169 .filterMode = CU_TR_FILTER_MODE_LINEAR,

170 .flags = CU_TRSF_READ_AS_INTEGER,

171 };

172

173 CUDA_RESOURCE_DESC res_desc = {

174 .resType = CU_RESOURCE_TYPE_PITCH2D,

175 .res.pitch2D.format = pixel_size == 1 ?

176 CU_AD_FORMAT_UNSIGNED_INT8 :

177 CU_AD_FORMAT_UNSIGNED_INT16,

178 .res.pitch2D.numChannels = channels,

179 .res.pitch2D.width = src_width,

180 .res.pitch2D.height = src_height,

181 .res.pitch2D.pitchInBytes = src_pitch,

182 .res.pitch2D.devPtr = (CUdeviceptr)src_dptr,

183 };

184

185 ret = CHECK_CU(cu->cuTexObjectCreate(&tex, &res_desc, &tex_desc, NULL));

186 if (ret < 0)

187 goto exit;

188

189 ret = CHECK_CU(cu->cuLaunchKernel(func,

190 DIV_UP(src_width, BLOCKX), DIV_UP(src_height, BLOCKY), 1,

191 BLOCKX, BLOCKY, 1, 0, s->cu_stream, args, NULL));

192 exit:

193 if (tex)

194 CHECK_CU(cu->cuTexObjectDestroy(tex));

195

196 return ret;

197 }

198

199 static int thumbnail(AVFilterContext *ctx, int *histogram, AVFrame *in)

200 {

201 AVHWFramesContext *in_frames_ctx = (AVHWFramesContext*)in->hw_frames_ctx->data;

202 ThumbnailCudaContext *s = ctx->priv;

203

204 switch (in_frames_ctx->sw_format) {

205 case AV_PIX_FMT_NV12:

206 thumbnail_kernel(ctx, s->cu_func_uchar, 1,

207 histogram, in->data[0], in->width, in->height, in->linesize[0], 1);

208 thumbnail_kernel(ctx, s->cu_func_uchar2, 2,

209 histogram + 256, in->data[1], in->width / 2, in->height / 2, in->linesize[1], 1);

210 break;

211 case AV_PIX_FMT_YUV420P:

212 thumbnail_kernel(ctx, s->cu_func_uchar, 1,

213 histogram, in->data[0], in->width, in->height, in->linesize[0], 1);

214 thumbnail_kernel(ctx, s->cu_func_uchar, 1,

215 histogram + 256, in->data[1], in->width / 2, in->height / 2, in->linesize[1], 1);

216 thumbnail_kernel(ctx, s->cu_func_uchar, 1,

217 histogram + 512, in->data[2], in->width / 2, in->height / 2, in->linesize[2], 1);

218 break;

219 case AV_PIX_FMT_YUV444P:

220 thumbnail_kernel(ctx, s->cu_func_uchar, 1,

221 histogram, in->data[0], in->width, in->height, in->linesize[0], 1);

222 thumbnail_kernel(ctx, s->cu_func_uchar, 1,

223 histogram + 256, in->data[1], in->width, in->height, in->linesize[1], 1);

224 thumbnail_kernel(ctx, s->cu_func_uchar, 1,

225 histogram + 512, in->data[2], in->width, in->height, in->linesize[2], 1);

226 break;

227 case AV_PIX_FMT_P010LE:

228 case AV_PIX_FMT_P016LE:

229 thumbnail_kernel(ctx, s->cu_func_ushort, 1,

230 histogram, in->data[0], in->width, in->height, in->linesize[0], 2);

231 thumbnail_kernel(ctx, s->cu_func_ushort2, 2,

232 histogram + 256, in->data[1], in->width / 2, in->height / 2, in->linesize[1], 2);

233 break;

234 case AV_PIX_FMT_YUV444P16:

235 thumbnail_kernel(ctx, s->cu_func_ushort2, 1,

236 histogram, in->data[0], in->width, in->height, in->linesize[0], 2);

237 thumbnail_kernel(ctx, s->cu_func_ushort2, 1,

238 histogram + 256, in->data[1], in->width, in->height, in->linesize[1], 2);

239 thumbnail_kernel(ctx, s->cu_func_ushort2, 1,

240 histogram + 512, in->data[2], in->width, in->height, in->linesize[2], 2);

241 break;

242 default:

243 return AVERROR_BUG;

244 }

245

246 return 0;

247 }

248

249 static int filter_frame(AVFilterLink *inlink, AVFrame *frame)

250 {

251 AVFilterContext *ctx = inlink->dst;

252 ThumbnailCudaContext *s = ctx->priv;

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

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

255 int *hist = s->frames[s->n].histogram;

256 AVHWFramesContext *hw_frames_ctx = (AVHWFramesContext*)s->hw_frames_ctx->data;

257 CUcontext dummy;

258 CUDA_MEMCPY2D cpy = { 0 };

259 int ret = 0;

260

261 // keep a reference of each frame

262 s->frames[s->n].buf = frame;

263

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

265 if (ret < 0)

266 return ret;

267

268 CHECK_CU(cu->cuMemsetD8Async(s->data, 0, HIST_SIZE * sizeof(int), s->cu_stream));

269

270 thumbnail(ctx, (int*)s->data, frame);

271

272 cpy.srcMemoryType = CU_MEMORYTYPE_DEVICE;

273 cpy.dstMemoryType = CU_MEMORYTYPE_HOST;

274 cpy.srcDevice = s->data;

275 cpy.dstHost = hist;

276 cpy.srcPitch = HIST_SIZE * sizeof(int);

277 cpy.dstPitch = HIST_SIZE * sizeof(int);

278 cpy.WidthInBytes = HIST_SIZE * sizeof(int);

279 cpy.Height = 1;

280

281 ret = CHECK_CU(cu->cuMemcpy2DAsync(&cpy, s->cu_stream));

282 if (ret < 0)

283 return ret;

284

285 if (hw_frames_ctx->sw_format == AV_PIX_FMT_NV12 || hw_frames_ctx->sw_format == AV_PIX_FMT_YUV420P ||

286 hw_frames_ctx->sw_format == AV_PIX_FMT_P010LE || hw_frames_ctx->sw_format == AV_PIX_FMT_P016LE)

287 {

288 int i;

289 for (i = 256; i < HIST_SIZE; i++)

290 hist[i] = 4 * hist[i];

291 }

292

293 ret = CHECK_CU(cu->cuCtxPopCurrent(&dummy));

294 if (ret < 0)

295 return ret;

296

297 // no selection until the buffer of N frames is filled up

298 s->n++;

299 if (s->n < s->n_frames)

300 return 0;

301

302 return ff_filter_frame(outlink, get_best_frame(ctx));

303 }

304

305 static av_cold void uninit(AVFilterContext *ctx)

306 {

307 ThumbnailCudaContext *s = ctx->priv;

308

309 if (s->hwctx) {

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

311

312 if (s->data) {

313 CHECK_CU(cu->cuMemFree(s->data));

314 s->data = 0;

315 }

316

317 if (s->cu_module) {

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

319 s->cu_module = NULL;

320 }

321 }

322

323 if (s->frames) {

324 for (int i = 0; i < s->n_frames && s->frames[i].buf; i++)

325 av_frame_free(&s->frames[i].buf);

326 av_freep(&s->frames);

327 }

328 }

329

330 static int request_frame(AVFilterLink *link)

331 {

332 AVFilterContext *ctx = link->src;

333 ThumbnailCudaContext *s = ctx->priv;

334 int ret = ff_request_frame(ctx->inputs[0]);

335

336 if (ret == AVERROR_EOF && s->n) {

337 ret = ff_filter_frame(link, get_best_frame(ctx));

338 if (ret < 0)

339 return ret;

340 ret = AVERROR_EOF;

341 }

342 if (ret < 0)

343 return ret;

344 return 0;

345 }

346

347 static int format_is_supported(enum AVPixelFormat fmt)

348 {

349 int i;

350

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

352 if (supported_formats[i] == fmt)

353 return 1;

354 return 0;

355 }

356

357 static int config_props(AVFilterLink *inlink)

358 {

359 AVFilterContext *ctx = inlink->dst;

360 ThumbnailCudaContext *s = ctx->priv;

361 AVHWFramesContext *hw_frames_ctx = (AVHWFramesContext*)inlink->hw_frames_ctx->data;

362 AVCUDADeviceContext *device_hwctx = hw_frames_ctx->device_ctx->hwctx;

363 CUcontext dummy, cuda_ctx = device_hwctx->cuda_ctx;

364 CudaFunctions *cu = device_hwctx->internal->cuda_dl;

365 int ret;

366

367 extern const unsigned char ff_vf_thumbnail_cuda_ptx_data[];

368 extern const unsigned int ff_vf_thumbnail_cuda_ptx_len;

369

370 s->hwctx = device_hwctx;

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

372

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

374 if (ret < 0)

375 return ret;

376

377 ret = ff_cuda_load_module(ctx, device_hwctx, &s->cu_module, ff_vf_thumbnail_cuda_ptx_data, ff_vf_thumbnail_cuda_ptx_len);

378 if (ret < 0)

379 return ret;

380

381 ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar, s->cu_module, "Thumbnail_uchar"));

382 if (ret < 0)

383 return ret;

384

385 ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar2, s->cu_module, "Thumbnail_uchar2"));

386 if (ret < 0)

387 return ret;

388

389 ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_ushort, s->cu_module, "Thumbnail_ushort"));

390 if (ret < 0)

391 return ret;

392

393 ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_ushort2, s->cu_module, "Thumbnail_ushort2"));

394 if (ret < 0)

395 return ret;

396

397 ret = CHECK_CU(cu->cuMemAlloc(&s->data, HIST_SIZE * sizeof(int)));

398 if (ret < 0)

399 return ret;

400

401 CHECK_CU(cu->cuCtxPopCurrent(&dummy));

402

403 s->hw_frames_ctx = ctx->inputs[0]->hw_frames_ctx;

404

405 ctx->outputs[0]->hw_frames_ctx = av_buffer_ref(s->hw_frames_ctx);

406 if (!ctx->outputs[0]->hw_frames_ctx)

407 return AVERROR(ENOMEM);

408

409 s->tb = inlink->time_base;

410

411 if (!format_is_supported(hw_frames_ctx->sw_format)) {

412 av_log(ctx, AV_LOG_ERROR, "Unsupported input format: %s\n", av_get_pix_fmt_name(hw_frames_ctx->sw_format));

413 return AVERROR(ENOSYS);

414 }

415

416 return 0;

417 }

418

419 static const AVFilterPad thumbnail_cuda_inputs[] = {

420 {

421 .name = "default",

422 .type = AVMEDIA_TYPE_VIDEO,

423 .config_props = config_props,

424 .filter_frame = filter_frame,

425 },

426 };

427

428 static const AVFilterPad thumbnail_cuda_outputs[] = {

429 {

430 .name = "default",

431 .type = AVMEDIA_TYPE_VIDEO,

432 .request_frame = request_frame,

433 },

434 };

435

436 const AVFilter ff_vf_thumbnail_cuda = {

437 .name = "thumbnail_cuda",

438 .description = NULL_IF_CONFIG_SMALL("Select the most representative frame in a given sequence of consecutive frames."),

439 .priv_size = sizeof(ThumbnailCudaContext),

440 .init = init,

441 .uninit = uninit,

442 FILTER_INPUTS(thumbnail_cuda_inputs),

443 FILTER_OUTPUTS(thumbnail_cuda_outputs),

444 FILTER_SINGLE_PIXFMT(AV_PIX_FMT_CUDA),

445 .priv_class = &thumbnail_cuda_class,

446 .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,

447 };

func

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

Definition: jacosubdec.c:68

AVHWDeviceContext::hwctx

void * hwctx

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

Definition: hwcontext.h:85

AV_PIX_FMT_CUDA

@ AV_PIX_FMT_CUDA

HW acceleration through CUDA.

Definition: pixfmt.h:260

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

thumbnail_cuda_options

static const AVOption thumbnail_cuda_options[]

Definition: vf_thumbnail_cuda.c:80

ThumbnailCudaContext::cu_module

CUmodule cu_module

Definition: vf_thumbnail_cuda.c:65

hwcontext_cuda_internal.h

ff_vf_thumbnail_cuda

const AVFilter ff_vf_thumbnail_cuda

Definition: vf_thumbnail_cuda.c:436

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

ff_filter_frame

int ff_filter_frame(AVFilterLink *link, AVFrame *frame)

Send a frame of data to the next filter.

Definition: avfilter.c:1018

filter_frame

static int filter_frame(AVFilterLink *inlink, AVFrame *frame)

Definition: vf_thumbnail_cuda.c:249

AVERROR_EOF

#define AVERROR_EOF

End of file.

Definition: error.h:57

AVBufferRef::data

uint8_t * data

The data buffer.

Definition: buffer.h:90

thumbnail_cuda_inputs

static const AVFilterPad thumbnail_cuda_inputs[]

Definition: vf_thumbnail_cuda.c:419

thumb_frame::histogram

int histogram[HIST_SIZE]

RGB color distribution histogram of the frame.

Definition: vf_thumbnail.c:39

get_best_frame

static AVFrame * get_best_frame(AVFilterContext *ctx)

Definition: vf_thumbnail_cuda.c:119

thumb_frame::buf

AVFrame * buf

cached frame

Definition: vf_thumbnail.c:38

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

AVFrame

This structure describes decoded (raw) audio or video data.

Definition: frame.h:344

pixdesc.h

AVFrame::pts

int64_t pts

Presentation timestamp in time_base units (time when frame should be shown to user).

Definition: frame.h:456

AVFrame::width

int width

Definition: frame.h:416

AVOption

AVOption.

Definition: opt.h:346

ff_request_frame

int ff_request_frame(AVFilterLink *link)

Request an input frame from the filter at the other end of the link.

Definition: avfilter.c:462

AV_LOG_VERBOSE

#define AV_LOG_VERBOSE

Detailed information.

Definition: log.h:196

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

AVFilterLink

A link between two filters.

Definition: avfilter.h:542

ThumbnailCudaContext::n

int n

current frame

Definition: vf_thumbnail_cuda.c:57

AVFrame::data

uint8_t * data[AV_NUM_DATA_POINTERS]

pointer to the picture/channel planes.

Definition: frame.h:365

thumb_frame

Definition: vf_thumbnail.c:37

config_props

static int config_props(AVFilterLink *inlink)

Definition: vf_thumbnail_cuda.c:357

dummy

int dummy

Definition: motion.c:66

frame_sum_square_err

static double frame_sum_square_err(const int *hist, const double *median)

Compute Sum-square deviation to estimate "closeness".

Definition: vf_thumbnail_cuda.c:107

supported_formats

static enum AVPixelFormat supported_formats[]

Definition: vf_thumbnail_cuda.c:41

AVFilterPad

A filter pad used for either input or output.

Definition: internal.h:33

request_frame

static int request_frame(AVFilterLink *link)

Definition: vf_thumbnail_cuda.c:330

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

#define s(width, name)

Definition: cbs_vp9.c:198

ThumbnailCudaContext::cu_func_ushort

CUfunction cu_func_ushort

Definition: vf_thumbnail_cuda.c:69

FLAGS

#define FLAGS

Definition: vf_thumbnail_cuda.c:78

AV_PIX_FMT_YUV444P16

#define AV_PIX_FMT_YUV444P16

Definition: pixfmt.h:491

av_q2d

static double av_q2d(AVRational a)

Convert an AVRational to a double.

Definition: rational.h:104

format_is_supported

static int format_is_supported(enum AVPixelFormat fmt)

Definition: vf_thumbnail_cuda.c:347

ctx

AVFormatContext * ctx

Definition: movenc.c:48

channels

Definition: aptx.h:31

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

Definition: filter_design.txt:23

frame

static AVFrame * frame

Definition: demux_decode.c:54

HIST_SIZE

#define HIST_SIZE

Definition: vf_thumbnail_cuda.c:36

ThumbnailCudaContext

Definition: vf_thumbnail_cuda.c:55

ThumbnailCudaContext::frames

struct thumb_frame * frames

the n_frames frames

Definition: vf_thumbnail_cuda.c:59

init

static av_cold int init(AVFilterContext *ctx)

Definition: vf_thumbnail_cuda.c:87

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

AVRational

Rational number (pair of numerator and denominator).

Definition: rational.h:58

double

Definition: af_crystalizer.c:131

AVFILTER_DEFINE_CLASS

AVFILTER_DEFINE_CLASS(thumbnail_cuda)

BLOCKY

#define BLOCKY

Definition: vf_thumbnail_cuda.c:39

ThumbnailCudaContext::hwctx

AVCUDADeviceContext * hwctx

Definition: vf_thumbnail_cuda.c:63

uninit

static av_cold void uninit(AVFilterContext *ctx)

Definition: vf_thumbnail_cuda.c:305

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

BLOCKX

#define BLOCKX

Definition: vf_thumbnail_cuda.c:38

ThumbnailCudaContext::cu_func_uchar

CUfunction cu_func_uchar

Definition: vf_thumbnail_cuda.c:67

ThumbnailCudaContext::tb

AVRational tb

copy of the input timebase to ease access

Definition: vf_thumbnail_cuda.c:60

AV_LOG_INFO

#define AV_LOG_INFO

Standard information.

Definition: log.h:191

internal.h

FILTER_SINGLE_PIXFMT

#define FILTER_SINGLE_PIXFMT(pix_fmt_)

Definition: internal.h:172

thumbnail

static int thumbnail(AVFilterContext *ctx, int *histogram, AVFrame *in)

Definition: vf_thumbnail_cuda.c:199

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

Definition: cbs_h2645.c:255

CHECK_CU

#define CHECK_CU(x)

Definition: vf_thumbnail_cuda.c:34

AVFilterPad::name

const char * name

Pad name.

Definition: internal.h:39

av_calloc

void * av_calloc(size_t nmemb, size_t size)

Definition: mem.c:262

AV_PIX_FMT_P016

#define AV_PIX_FMT_P016

Definition: pixfmt.h:530

AVFilter

Filter definition.

Definition: avfilter.h:166

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

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

AVHWFramesContext::device_ctx

AVHWDeviceContext * device_ctx

The parent AVHWDeviceContext.

Definition: hwcontext.h:134

cuda_check.h

AV_PIX_FMT_P016LE

@ AV_PIX_FMT_P016LE

like NV12, with 16bpp per component, little-endian

Definition: pixfmt.h:323

ThumbnailCudaContext::data

CUdeviceptr data

Definition: vf_thumbnail_cuda.c:73

AVFrame::hw_frames_ctx

AVBufferRef * hw_frames_ctx

For hwaccel-format frames, this should be a reference to the AVHWFramesContext describing the frame.

Definition: frame.h:695

AVFrame::height

int height

Definition: frame.h:416

AV_OPT_TYPE_INT

@ AV_OPT_TYPE_INT

Definition: opt.h:235

avfilter.h

thumbnail_cuda_outputs

static const AVFilterPad thumbnail_cuda_outputs[]

Definition: vf_thumbnail_cuda.c:428

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

ThumbnailCudaContext::cu_func_ushort2

CUfunction cu_func_ushort2

Definition: vf_thumbnail_cuda.c:70

AV_PIX_FMT_P010

#define AV_PIX_FMT_P010

Definition: pixfmt.h:528

AVMEDIA_TYPE_VIDEO

@ AVMEDIA_TYPE_VIDEO

Definition: avutil.h:201

AVBufferRef

A reference to a data buffer.

Definition: buffer.h:82

AV_PIX_FMT_P010LE

@ AV_PIX_FMT_P010LE

like NV12, with 10bpp per component, data in the high bits, zeros in the low bits,...

Definition: pixfmt.h:307

ThumbnailCudaContext::cu_stream

CUstream cu_stream

Definition: vf_thumbnail_cuda.c:71

thumbnail_kernel

static int thumbnail_kernel(AVFilterContext *ctx, CUfunction func, int channels, int *histogram, uint8_t *src_dptr, int src_width, int src_height, int src_pitch, int pixel_size)

Definition: vf_thumbnail_cuda.c:159

FILTER_OUTPUTS

#define FILTER_OUTPUTS(array)

Definition: internal.h:183

av_freep

#define av_freep(p)

Definition: tableprint_vlc.h:34

ThumbnailCudaContext::n_frames

int n_frames

number of frames for analysis

Definition: vf_thumbnail_cuda.c:58

ThumbnailCudaContext::hw_frames_ctx