FFmpeg: libavfilter/vf_nlmeans_opencl.c Source File

FFmpeg

[フレーム]

libavfilter

vf_nlmeans_opencl.c

Go to the documentation of this file.

1 /*

2 * This file is part of FFmpeg.

3 *

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

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

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

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

8 *

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

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

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

12 * Lesser General Public License for more details.

13 *

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

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

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

17 */

18 #include <float.h>

20 #include "libavutil/avassert.h"

21 #include "libavutil/common.h"

22 #include "libavutil/imgutils.h"

23 #include "libavutil/mem.h"

24 #include "libavutil/opt.h"

25 #include "libavutil/pixdesc.h"

27 #include "avfilter.h"

28 #include "internal.h"

29 #include "opencl.h"

30 #include "opencl_source.h"

31 #include "video.h"

33 // TODO:

34 // the integral image may overflow 32bit, consider using 64bit

36 static const enum AVPixelFormat supported_formats[] = {

37 AV_PIX_FMT_YUV420P,

38 AV_PIX_FMT_YUV444P,

39 AV_PIX_FMT_GBRP,

40 };

42 static int is_format_supported(enum AVPixelFormat fmt)

43 {

44 int i;

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

47 if (supported_formats[i] == fmt)

48 return 1;

49 return 0;

50 }

52 typedef struct NLMeansOpenCLContext {

53 OpenCLFilterContext ocf;

54 int initialised;

55 cl_kernel vert_kernel;

56 cl_kernel horiz_kernel;

57 cl_kernel accum_kernel;

58 cl_kernel average_kernel;

59 cl_mem integral_img;

60 cl_mem weight;

61 cl_mem sum;

62 cl_mem overflow; // overflow in integral image?

63 double sigma;

64 float h;

65 int chroma_w;

66 int chroma_h;

67 int patch_size;

68 int patch_size_uv;

69 int research_size;

70 int research_size_uv;

71 cl_command_queue command_queue;

72 } NLMeansOpenCLContext;

74 static int nlmeans_opencl_init(AVFilterContext *avctx, int width, int height)

75 {

76 NLMeansOpenCLContext *ctx = avctx->priv;

77 cl_int cle;

78 int err;

79 int weight_buf_size = width * height * sizeof(float);

81 ctx->h = ctx->sigma * 10;

82 if (!(ctx->research_size & 1)) {

83 ctx->research_size |= 1;

84 av_log(avctx, AV_LOG_WARNING,

85 "research_size should be odd, set to %d",

86 ctx->research_size);

87 }

89 if (!(ctx->patch_size & 1)) {

90 ctx->patch_size |= 1;

91 av_log(avctx, AV_LOG_WARNING,

92 "patch_size should be odd, set to %d",

93 ctx->patch_size);

94 }

96 if (!ctx->research_size_uv)

97 ctx->research_size_uv = ctx->research_size;

98 if (!ctx->patch_size_uv)

99 ctx->patch_size_uv = ctx->patch_size;

100

101 err = ff_opencl_filter_load_program(avctx, &ff_opencl_source_nlmeans, 1);

102 if (err < 0)

103 goto fail;

104

105 ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,

106 ctx->ocf.hwctx->device_id,

107 0, &cle);

108 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL "

109 "command queue %d.\n", cle);

110

111 ctx->vert_kernel = clCreateKernel(ctx->ocf.program,

112 "vert_sum", &cle);

113 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create "

114 "vert_sum kernel %d.\n", cle);

115

116 ctx->horiz_kernel = clCreateKernel(ctx->ocf.program,

117 "horiz_sum", &cle);

118 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create "

119 "horiz_sum kernel %d.\n", cle);

120

121 ctx->accum_kernel = clCreateKernel(ctx->ocf.program,

122 "weight_accum", &cle);

123 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create "

124 "accum kernel %d.\n", cle);

125

126 ctx->average_kernel = clCreateKernel(ctx->ocf.program,

127 "average", &cle);

128 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create "

129 "average kernel %d.\n", cle);

130

131 ctx->integral_img = clCreateBuffer(ctx->ocf.hwctx->context, 0,

132 4 * width * height * sizeof(cl_int),

133 NULL, &cle);

134 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create "

135 "integral image %d.\n", cle);

136

137 ctx->weight = clCreateBuffer(ctx->ocf.hwctx->context, 0,

138 weight_buf_size, NULL, &cle);

139 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create "

140 "weight buffer %d.\n", cle);

141

142 ctx->sum = clCreateBuffer(ctx->ocf.hwctx->context, 0,

143 weight_buf_size, NULL, &cle);

144 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create "

145 "sum buffer %d.\n", cle);

146

147 ctx->overflow = clCreateBuffer(ctx->ocf.hwctx->context, 0,

148 sizeof(cl_int), NULL, &cle);

149 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create "

150 "overflow buffer %d.\n", cle);

151

152 ctx->initialised = 1;

153 return 0;

154

155 fail:

156 CL_RELEASE_KERNEL(ctx->vert_kernel);

157 CL_RELEASE_KERNEL(ctx->horiz_kernel);

158 CL_RELEASE_KERNEL(ctx->accum_kernel);

159 CL_RELEASE_KERNEL(ctx->average_kernel);

160

161 CL_RELEASE_MEMORY(ctx->integral_img);

162 CL_RELEASE_MEMORY(ctx->weight);

163 CL_RELEASE_MEMORY(ctx->sum);

164 CL_RELEASE_MEMORY(ctx->overflow);

165

166 CL_RELEASE_QUEUE(ctx->command_queue);

167 return err;

168 }

169

170 static int nlmeans_plane(AVFilterContext *avctx, cl_mem dst, cl_mem src,

171 cl_int width, cl_int height, cl_int p, cl_int r)

172 {

173 NLMeansOpenCLContext *ctx = avctx->priv;

174 const float zero = 0.0f;

175 const size_t worksize1[] = {height};

176 const size_t worksize2[] = {width};

177 const size_t worksize3[2] = {width, height};

178 int i, dx, dy, err = 0, weight_buf_size;

179 cl_int cle;

180 int nb_pixel, *tmp = NULL, idx = 0;

181 cl_int *dxdy = NULL;

182

183 weight_buf_size = width * height * sizeof(float);

184 cle = clEnqueueFillBuffer(ctx->command_queue, ctx->weight,

185 &zero, sizeof(float), 0, weight_buf_size,

186 0, NULL, NULL);

187 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to fill weight buffer: %d.\n",

188 cle);

189 cle = clEnqueueFillBuffer(ctx->command_queue, ctx->sum,

190 &zero, sizeof(float), 0, weight_buf_size,

191 0, NULL, NULL);

192 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to fill sum buffer: %d.\n",

193 cle);

194

195 nb_pixel = (2 * r + 1) * (2 * r + 1) - 1;

196 dxdy = av_malloc(nb_pixel * 2 * sizeof(cl_int));

197 tmp = av_malloc(nb_pixel * 2 * sizeof(int));

198

199 if (!dxdy || !tmp)

200 goto fail;

201

202 for (dx = -r; dx <= r; dx++) {

203 for (dy = -r; dy <= r; dy++) {

204 if (dx || dy) {

205 tmp[idx++] = dx;

206 tmp[idx++] = dy;

207 }

208 }

209 }

210 // repack dx/dy seperately, as we want to do four pairs of dx/dy in a batch

211 for (i = 0; i < nb_pixel / 4; i++) {

212 dxdy[i * 8] = tmp[i * 8]; // dx0

213 dxdy[i * 8 + 1] = tmp[i * 8 + 2]; // dx1

214 dxdy[i * 8 + 2] = tmp[i * 8 + 4]; // dx2

215 dxdy[i * 8 + 3] = tmp[i * 8 + 6]; // dx3

216 dxdy[i * 8 + 4] = tmp[i * 8 + 1]; // dy0

217 dxdy[i * 8 + 5] = tmp[i * 8 + 3]; // dy1

218 dxdy[i * 8 + 6] = tmp[i * 8 + 5]; // dy2

219 dxdy[i * 8 + 7] = tmp[i * 8 + 7]; // dy3

220 }

221 av_freep(&tmp);

222

223 for (i = 0; i < nb_pixel / 4; i++) {

224 cl_int *dx_cur = dxdy + 8 * i;

225 cl_int *dy_cur = dxdy + 8 * i + 4;

226

227 // horizontal pass

228 // integral(x,y) = sum([u(v,y) - u(v+dx,y+dy)]^2) for v in [0, x]

229 CL_SET_KERNEL_ARG(ctx->horiz_kernel, 0, cl_mem, &ctx->integral_img);

230 CL_SET_KERNEL_ARG(ctx->horiz_kernel, 1, cl_mem, &src);

231 CL_SET_KERNEL_ARG(ctx->horiz_kernel, 2, cl_int, &width);

232 CL_SET_KERNEL_ARG(ctx->horiz_kernel, 3, cl_int, &height);

233 CL_SET_KERNEL_ARG(ctx->horiz_kernel, 4, cl_int4, dx_cur);

234 CL_SET_KERNEL_ARG(ctx->horiz_kernel, 5, cl_int4, dy_cur);

235 cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->horiz_kernel, 1,

236 NULL, worksize1, NULL, 0, NULL, NULL);

237 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue horiz_kernel: %d.\n",

238 cle);

239 // vertical pass

240 // integral(x, y) = sum(integral(x, v)) for v in [0, y]

241 CL_SET_KERNEL_ARG(ctx->vert_kernel, 0, cl_mem, &ctx->integral_img);

242 CL_SET_KERNEL_ARG(ctx->vert_kernel, 1, cl_mem, &ctx->overflow);

243 CL_SET_KERNEL_ARG(ctx->vert_kernel, 2, cl_int, &width);

244 CL_SET_KERNEL_ARG(ctx->vert_kernel, 3, cl_int, &height);

245 cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->vert_kernel,

246 1, NULL, worksize2, NULL, 0, NULL, NULL);

247 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue vert_kernel: %d.\n",

248 cle);

249

250 // accumulate weights

251 CL_SET_KERNEL_ARG(ctx->accum_kernel, 0, cl_mem, &ctx->sum);

252 CL_SET_KERNEL_ARG(ctx->accum_kernel, 1, cl_mem, &ctx->weight);

253 CL_SET_KERNEL_ARG(ctx->accum_kernel, 2, cl_mem, &ctx->integral_img);

254 CL_SET_KERNEL_ARG(ctx->accum_kernel, 3, cl_mem, &src);

255 CL_SET_KERNEL_ARG(ctx->accum_kernel, 4, cl_int, &width);

256 CL_SET_KERNEL_ARG(ctx->accum_kernel, 5, cl_int, &height);

257 CL_SET_KERNEL_ARG(ctx->accum_kernel, 6, cl_int, &p);

258 CL_SET_KERNEL_ARG(ctx->accum_kernel, 7, cl_float, &ctx->h);

259 CL_SET_KERNEL_ARG(ctx->accum_kernel, 8, cl_int4, dx_cur);

260 CL_SET_KERNEL_ARG(ctx->accum_kernel, 9, cl_int4, dy_cur);

261 cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->accum_kernel,

262 2, NULL, worksize3, NULL, 0, NULL, NULL);

263 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue kernel: %d.\n", cle);

264 }

265 av_freep(&dxdy);

266

267 // average

268 CL_SET_KERNEL_ARG(ctx->average_kernel, 0, cl_mem, &dst);

269 CL_SET_KERNEL_ARG(ctx->average_kernel, 1, cl_mem, &src);

270 CL_SET_KERNEL_ARG(ctx->average_kernel, 2, cl_mem, &ctx->sum);

271 CL_SET_KERNEL_ARG(ctx->average_kernel, 3, cl_mem, &ctx->weight);

272 cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->average_kernel, 2,

273 NULL, worksize3, NULL, 0, NULL, NULL);

274 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue average kernel: %d.\n",

275 cle);

276 cle = clFlush(ctx->command_queue);

277 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to flush command queue: %d.\n", cle);

278 fail:

279 if (tmp)

280 av_freep(&tmp);

281 if (dxdy)

282 av_freep(&dxdy);

283 return err;

284 }

285

286 static int nlmeans_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)

287 {

288 AVFilterContext *avctx = inlink->dst;

289 AVFilterLink *outlink = avctx->outputs[0];

290 NLMeansOpenCLContext *ctx = avctx->priv;

291 AVFrame *output = NULL;

292 AVHWFramesContext *input_frames_ctx;

293 const AVPixFmtDescriptor *desc;

294 enum AVPixelFormat in_format;

295 cl_mem src, dst;

296 const cl_int zero = 0;

297 int w, h, err, cle, overflow, p, patch, research;

298

299 av_log(ctx, AV_LOG_DEBUG, "Filter input: %s, %ux%u (%"PRId64").\n",

300 av_get_pix_fmt_name(input->format),

301 input->width, input->height, input->pts);

302

303 if (!input->hw_frames_ctx)

304 return AVERROR(EINVAL);

305 input_frames_ctx = (AVHWFramesContext*)input->hw_frames_ctx->data;

306 in_format = input_frames_ctx->sw_format;

307

308 output = ff_get_video_buffer(outlink, outlink->w, outlink->h);

309 if (!output) {

310 err = AVERROR(ENOMEM);

311 goto fail;

312 }

313

314 err = av_frame_copy_props(output, input);

315 if (err < 0)

316 goto fail;

317

318 if (!ctx->initialised) {

319 desc = av_pix_fmt_desc_get(in_format);

320 if (!is_format_supported(in_format)) {

321 err = AVERROR(EINVAL);

322 av_log(avctx, AV_LOG_ERROR, "input format %s not supported\n",

323 av_get_pix_fmt_name(in_format));

324 goto fail;

325 }

326 ctx->chroma_w = AV_CEIL_RSHIFT(inlink->w, desc->log2_chroma_w);

327 ctx->chroma_h = AV_CEIL_RSHIFT(inlink->h, desc->log2_chroma_h);

328

329 err = nlmeans_opencl_init(avctx, inlink->w, inlink->h);

330 if (err < 0)

331 goto fail;

332 }

333

334 cle = clEnqueueWriteBuffer(ctx->command_queue, ctx->overflow, CL_FALSE,

335 0, sizeof(cl_int), &zero, 0, NULL, NULL);

336 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to initialize overflow"

337 "detection buffer %d.\n", cle);

338

339 for (p = 0; p < FF_ARRAY_ELEMS(output->data); p++) {

340 src = (cl_mem) input->data[p];

341 dst = (cl_mem) output->data[p];

342

343 if (!dst)

344 break;

345 av_assert0(src);

346 w = p ? ctx->chroma_w : inlink->w;

347 h = p ? ctx->chroma_h : inlink->h;

348 patch = (p ? ctx->patch_size_uv : ctx->patch_size) / 2;

349 research = (p ? ctx->research_size_uv : ctx->research_size) / 2;

350 err = nlmeans_plane(avctx, dst, src, w, h, patch, research);

351 if (err < 0)

352 goto fail;

353 }

354 // overflow occurred?

355 cle = clEnqueueReadBuffer(ctx->command_queue, ctx->overflow, CL_FALSE,

356 0, sizeof(cl_int), &overflow, 0, NULL, NULL);

357 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to read overflow: %d.\n", cle);

358

359 cle = clFinish(ctx->command_queue);

360 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish kernel: %d.\n", cle);

361

362 if (overflow > 0)

363 av_log(avctx, AV_LOG_ERROR, "integral image overflow %d\n", overflow);

364

365 av_frame_free(&input);

366

367 av_log(ctx, AV_LOG_DEBUG, "Filter output: %s, %ux%u (%"PRId64").\n",

368 av_get_pix_fmt_name(output->format),

369 output->width, output->height, output->pts);

370

371 return ff_filter_frame(outlink, output);

372

373 fail:

374 clFinish(ctx->command_queue);

375 av_frame_free(&input);

376 av_frame_free(&output);

377 return err;

378 }

379

380 static av_cold void nlmeans_opencl_uninit(AVFilterContext *avctx)

381 {

382 NLMeansOpenCLContext *ctx = avctx->priv;

383 cl_int cle;

384

385 CL_RELEASE_KERNEL(ctx->vert_kernel);

386 CL_RELEASE_KERNEL(ctx->horiz_kernel);

387 CL_RELEASE_KERNEL(ctx->accum_kernel);

388 CL_RELEASE_KERNEL(ctx->average_kernel);

389

390 CL_RELEASE_MEMORY(ctx->integral_img);

391 CL_RELEASE_MEMORY(ctx->weight);

392 CL_RELEASE_MEMORY(ctx->sum);

393 CL_RELEASE_MEMORY(ctx->overflow);

394

395 CL_RELEASE_QUEUE(ctx->command_queue);

396

397 ff_opencl_filter_uninit(avctx);

398 }

399

400 #define OFFSET(x) offsetof(NLMeansOpenCLContext, x)

401 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)

402 static const AVOption nlmeans_opencl_options[] = {

403 { "s", "denoising strength", OFFSET(sigma), AV_OPT_TYPE_DOUBLE, { .dbl = 1.0 }, 1.0, 30.0, FLAGS },

404 { "p", "patch size", OFFSET(patch_size), AV_OPT_TYPE_INT, { .i64 = 2*3+1 }, 0, 99, FLAGS },

405 { "pc", "patch size for chroma planes", OFFSET(patch_size_uv), AV_OPT_TYPE_INT, { .i64 = 0 }, 0, 99, FLAGS },

406 { "r", "research window", OFFSET(research_size), AV_OPT_TYPE_INT, { .i64 = 7*2+1 }, 0, 99, FLAGS },

407 { "rc", "research window for chroma planes", OFFSET(research_size_uv), AV_OPT_TYPE_INT, { .i64 = 0 }, 0, 99, FLAGS },

408 { NULL }

409 };

410

411 AVFILTER_DEFINE_CLASS(nlmeans_opencl);

412

413 static const AVFilterPad nlmeans_opencl_inputs[] = {

414 {

415 .name = "default",

416 .type = AVMEDIA_TYPE_VIDEO,

417 .filter_frame = &nlmeans_opencl_filter_frame,

418 .config_props = &ff_opencl_filter_config_input,

419 },

420 };

421

422 static const AVFilterPad nlmeans_opencl_outputs[] = {

423 {

424 .name = "default",

425 .type = AVMEDIA_TYPE_VIDEO,

426 .config_props = &ff_opencl_filter_config_output,

427 },

428 };

429

430 const AVFilter ff_vf_nlmeans_opencl = {

431 .name = "nlmeans_opencl",

432 .description = NULL_IF_CONFIG_SMALL("Non-local means denoiser through OpenCL"),

433 .priv_size = sizeof(NLMeansOpenCLContext),

434 .priv_class = &nlmeans_opencl_class,

435 .init = &ff_opencl_filter_init,

436 .uninit = &nlmeans_opencl_uninit,

437 FILTER_INPUTS(nlmeans_opencl_inputs),

438 FILTER_OUTPUTS(nlmeans_opencl_outputs),

439 FILTER_SINGLE_PIXFMT(AV_PIX_FMT_OPENCL),

440 .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,

441 };

ff_get_video_buffer

AVFrame * ff_get_video_buffer(AVFilterLink *link, int w, int h)

Request a picture buffer with a specific set of permissions.

Definition: video.c:101

NLMeansOpenCLContext::horiz_kernel

cl_kernel horiz_kernel

Definition: vf_nlmeans_opencl.c:56

AV_LOG_WARNING

#define AV_LOG_WARNING

Something somehow does not look correct.

Definition: log.h:186

nlmeans_opencl_options

static const AVOption nlmeans_opencl_options[]

Definition: vf_nlmeans_opencl.c:402

AVPixelFormat

Pixel format.

Definition: pixfmt.h:64

NLMeansOpenCLContext::research_size_uv

int research_size_uv

Definition: vf_nlmeans_opencl.c:70

const char * r

Definition: vf_curves.c:126

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

NLMeansOpenCLContext::chroma_h

int chroma_h

Definition: vf_nlmeans_opencl.c:66

nlmeans_opencl_init

static int nlmeans_opencl_init(AVFilterContext *avctx, int width, int height)

Definition: vf_nlmeans_opencl.c:74

CL_SET_KERNEL_ARG

#define CL_SET_KERNEL_ARG(kernel, arg_num, type, arg)

set argument to specific Kernel.

Definition: opencl.h:61

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

ff_filter_frame

int ff_filter_frame(AVFilterLink *link, AVFrame *frame)

Send a frame of data to the next filter.

Definition: avfilter.c:969

AVFILTER_DEFINE_CLASS

AVFILTER_DEFINE_CLASS(nlmeans_opencl)

NLMeansOpenCLContext

Definition: vf_nlmeans_opencl.c:52

av_pix_fmt_desc_get

const AVPixFmtDescriptor * av_pix_fmt_desc_get(enum AVPixelFormat pix_fmt)

Definition: pixdesc.c:2888

output

filter_frame For filters that do not use the this method is called when a frame is pushed to the filter s input It can be called at any time except in a reentrant way If the input frame is enough to produce output

Definition: filter_design.txt:225

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

test::height

int height

Definition: vc1dsp.c:39

nlmeans_opencl_outputs

static const AVFilterPad nlmeans_opencl_outputs[]

Definition: vf_nlmeans_opencl.c:422

AVFrame

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

Definition: frame.h:330

tmp

static uint8_t tmp[11]

Definition: aes_ctr.c:28

pixdesc.h

uint8_t w

Definition: llviddspenc.c:38

opencl.h

AVOption

AVOption.

Definition: opt.h:251

float.h

ff_opencl_filter_load_program

int ff_opencl_filter_load_program(AVFilterContext *avctx, const char **program_source_array, int nb_strings)

Load a new OpenCL program from strings in memory.

Definition: opencl.c:157

is_format_supported

static int is_format_supported(enum AVPixelFormat fmt)

Definition: vf_nlmeans_opencl.c:42

AVFilter::name

const char * name

Filter name.

Definition: avfilter.h:165

video.h

AVFilterLink

A link between two filters.

Definition: avfilter.h:522

ff_vf_nlmeans_opencl

const AVFilter ff_vf_nlmeans_opencl

Definition: vf_nlmeans_opencl.c:430

av_malloc

#define av_malloc(s)

Definition: tableprint_vlc.h:30

NLMeansOpenCLContext::patch_size

int patch_size

Definition: vf_nlmeans_opencl.c:67

AVFilterContext::priv

void * priv

private data for use by the filter

Definition: avfilter.h:407

fail

#define fail()

Definition: checkasm.h:134

FLAGS

#define FLAGS

Definition: vf_nlmeans_opencl.c:401

NLMeansOpenCLContext::patch_size_uv

int patch_size_uv

Definition: vf_nlmeans_opencl.c:68

ff_opencl_filter_config_output

int ff_opencl_filter_config_output(AVFilterLink *outlink)

Create a suitable hardware frames context for the output.

Definition: opencl.c:82

AVFilterPad

A filter pad used for either input or output.

Definition: internal.h:49

avassert.h

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

float

Definition: af_crystalizer.c:122

width

#define width

AV_CEIL_RSHIFT

#define AV_CEIL_RSHIFT(a, b)

Definition: common.h:50

AV_OPT_TYPE_DOUBLE

@ AV_OPT_TYPE_DOUBLE

Definition: opt.h:227

NLMeansOpenCLContext::h

float h

Definition: vf_nlmeans_opencl.c:64

init

int(* init)(AVBSFContext *ctx)

Definition: dts2pts_bsf.c:365

av_assert0

#define av_assert0(cond)

assert() equivalent, that is always enabled.

Definition: avassert.h:37

AV_LOG_DEBUG

#define AV_LOG_DEBUG

Stuff which is only useful for libav* developers.

Definition: log.h:201

NLMeansOpenCLContext::integral_img

cl_mem integral_img

Definition: vf_nlmeans_opencl.c:59

ctx

AVFormatContext * ctx

Definition: movenc.c:48

AV_PIX_FMT_YUV420P

@ AV_PIX_FMT_YUV420P

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

Definition: pixfmt.h:66

FILTER_INPUTS

#define FILTER_INPUTS(array)

Definition: internal.h:194

NLMeansOpenCLContext::weight

cl_mem weight

Definition: vf_nlmeans_opencl.c:60

CL_RELEASE_KERNEL

#define CL_RELEASE_KERNEL(k)

release an OpenCL Kernel

Definition: opencl.h:101

NULL

#define NULL

Definition: coverity.c:32

AVHWFramesContext::sw_format

enum AVPixelFormat sw_format

The pixel format identifying the actual data layout of the hardware frames.

Definition: hwcontext.h:222

av_frame_copy_props

int av_frame_copy_props(AVFrame *dst, const AVFrame *src)

Copy only "metadata" fields from src to dst.

Definition: frame.c:594

AV_PIX_FMT_OPENCL

@ AV_PIX_FMT_OPENCL

Hardware surfaces for OpenCL.

Definition: pixfmt.h:355

NLMeansOpenCLContext::vert_kernel

cl_kernel vert_kernel

Definition: vf_nlmeans_opencl.c:55

NLMeansOpenCLContext::initialised

int initialised

Definition: vf_nlmeans_opencl.c:54

NLMeansOpenCLContext::sum

cl_mem sum

Definition: vf_nlmeans_opencl.c:61

test::width

int width

Definition: vc1dsp.c:38

NULL_IF_CONFIG_SMALL

#define NULL_IF_CONFIG_SMALL(x)

Return NULL if CONFIG_SMALL is true, otherwise the argument without modification.

Definition: internal.h:115

CL_RELEASE_MEMORY

#define CL_RELEASE_MEMORY(m)

release an OpenCL Memory Object

Definition: opencl.h:114

nlmeans_opencl_inputs

static const AVFilterPad nlmeans_opencl_inputs[]

Definition: vf_nlmeans_opencl.c:413

NLMeansOpenCLContext::accum_kernel

cl_kernel accum_kernel

Definition: vf_nlmeans_opencl.c:57

NLMeansOpenCLContext::average_kernel

cl_kernel average_kernel

Definition: vf_nlmeans_opencl.c:58

NLMeansOpenCLContext::overflow

cl_mem overflow

Definition: vf_nlmeans_opencl.c:62

NLMeansOpenCLContext::sigma

double sigma

Definition: vf_nlmeans_opencl.c:63

height

#define height

ff_opencl_source_nlmeans

const char * ff_opencl_source_nlmeans

nlmeans_opencl_uninit

static av_cold void nlmeans_opencl_uninit(AVFilterContext *avctx)

Definition: vf_nlmeans_opencl.c:380

opencl_source.h

input

and forward the test the status of outputs and forward it to the corresponding return FFERROR_NOT_READY If the filters stores internally one or a few frame for some input

Definition: filter_design.txt:172

ff_opencl_filter_config_input

int ff_opencl_filter_config_input(AVFilterLink *inlink)

Check that the input link contains a suitable hardware frames context and extract the device from it.

Definition: opencl.c:46

internal.h

FILTER_SINGLE_PIXFMT

#define FILTER_SINGLE_PIXFMT(pix_fmt_)

Definition: internal.h:184

NLMeansOpenCLContext::command_queue

cl_command_queue command_queue

Definition: vf_nlmeans_opencl.c:71