FFmpeg: libavfilter/vf_convolution_opencl.c Source File

FFmpeg

[フレーム]

libavfilter

vf_convolution_opencl.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 "config_components.h"

23 #include "libavutil/avassert.h"

24 #include "libavutil/common.h"

25 #include "libavutil/imgutils.h"

26 #include "libavutil/mem.h"

27 #include "libavutil/opt.h"

28 #include "libavutil/pixdesc.h"

29 #include "libavutil/avstring.h"

32 #include "avfilter.h"

33 #include "filters.h"

34 #include "opencl.h"

35 #include "opencl_source.h"

36 #include "video.h"

38 typedef struct ConvolutionOpenCLContext {

39 OpenCLFilterContext ocf;

41 int initialised;

42 cl_kernel kernel;

43 cl_command_queue command_queue;

45 char *matrix_str[4];

47 cl_mem matrix[4];

48 cl_int matrix_sizes[4];

49 cl_int dims[4];

50 cl_float rdivs[4];

51 cl_float biases[4];

53 cl_int planes;

54 cl_float scale;

55 cl_float delta;

57 } ConvolutionOpenCLContext;

59 static int convolution_opencl_init(AVFilterContext *avctx)

60 {

61 ConvolutionOpenCLContext *ctx = avctx->priv;

62 const char *kernel_name;

63 cl_int cle;

64 int err;

66 err = ff_opencl_filter_load_program(avctx, &ff_source_convolution_cl, 1);

67 if (err < 0)

68 goto fail;

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

71 ctx->ocf.hwctx->device_id,

72 0, &cle);

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

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

76 if (!strcmp(avctx->filter->name, "convolution_opencl")) {

77 kernel_name = "convolution_global";

78 } else if (!strcmp(avctx->filter->name, "sobel_opencl")) {

79 kernel_name = "sobel_global";

80 } else if (!strcmp(avctx->filter->name, "prewitt_opencl")){

81 kernel_name = "prewitt_global";

82 } else if (!strcmp(avctx->filter->name, "roberts_opencl")){

83 kernel_name = "roberts_global";

84 } else {

85 av_assert0(0);

86 }

87 ctx->kernel = clCreateKernel(ctx->ocf.program, kernel_name, &cle);

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

89 "kernel %d.\n", cle);

91 ctx->initialised = 1;

92 return 0;

94 fail:

95 if (ctx->command_queue)

96 clReleaseCommandQueue(ctx->command_queue);

97 if (ctx->kernel)

98 clReleaseKernel(ctx->kernel);

99 return err;

100 }

101

102

103

104 static int convolution_opencl_make_filter_params(AVFilterContext *avctx)

105 {

106 ConvolutionOpenCLContext *ctx = avctx->priv;

107 float *matrix = NULL;

108 size_t matrix_bytes;

109 cl_mem buffer;

110 cl_int cle;

111 int i, j;

112 int sscanf_err;

113 char *p, *arg, *saveptr = NULL;

114 float input_matrix[4][49];

115

116 for (i = 0; i < 4; i++) {

117 ctx->biases[i] = ctx->biases[i] / 255.0;

118 }

119

120 for (i = 0; i < 4; i++) {

121 p = ctx->matrix_str[i];

122 while (ctx->matrix_sizes[i] < 49) {

123 arg = av_strtok(p, " ", &saveptr);

124 if (!arg) {

125 break;

126 }

127 p = NULL;

128 sscanf_err = sscanf(arg, "%f", &input_matrix[i][ctx->matrix_sizes[i]]);

129 if (sscanf_err != 1) {

130 av_log(ctx, AV_LOG_ERROR, "Matrix is sequence of 9, 25 or 49 signed numbers\n");

131 return AVERROR(EINVAL);

132 }

133 ctx->matrix_sizes[i]++;

134 }

135 if (ctx->matrix_sizes[i] == 9) {

136 ctx->dims[i] = 3;

137 } else if (ctx->matrix_sizes[i] == 25) {

138 ctx->dims[i] = 5;

139 } else if (ctx->matrix_sizes[i] == 49) {

140 ctx->dims[i] = 7;

141 } else {

142 av_log(ctx, AV_LOG_ERROR, "Invalid matrix size:%d\n", ctx->matrix_sizes[i]);

143 return AVERROR(EINVAL);

144 }

145

146 }

147

148 for (j = 0; j < 4; j++) {

149 matrix_bytes = sizeof(float)*ctx->matrix_sizes[j];

150 matrix = av_malloc(matrix_bytes);

151 if (!matrix) {

152 av_freep(&matrix);

153 return AVERROR(ENOMEM);

154 }

155

156 for (i = 0; i < ctx->matrix_sizes[j]; i++)

157 matrix[i] = input_matrix[j][i];

158

159 buffer = clCreateBuffer(ctx->ocf.hwctx->context,

160 CL_MEM_READ_ONLY |

161 CL_MEM_COPY_HOST_PTR |

162 CL_MEM_HOST_NO_ACCESS,

163 matrix_bytes, matrix, &cle);

164 if (!buffer) {

165 av_log(avctx, AV_LOG_ERROR, "Failed to create matrix buffer: "

166 "%d.\n", cle);

167 av_freep(&matrix);

168 return AVERROR(EIO);

169 }

170 ctx->matrix[j] = buffer;

171 av_freep(&matrix);

172 }

173

174 return 0;

175 }

176

177 static int convolution_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)

178 {

179 AVFilterContext *avctx = inlink->dst;

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

181 ConvolutionOpenCLContext *ctx = avctx->priv;

182 AVFrame *output = NULL;

183 cl_int cle;

184 size_t global_work[2];

185 cl_mem src, dst;

186 int err, p;

187 size_t origin[3] = {0, 0, 0};

188 size_t region[3] = {0, 0, 1};

189

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

191 av_get_pix_fmt_name(input->format),

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

193

194 if (!input->hw_frames_ctx)

195 return AVERROR(EINVAL);

196

197 if (!ctx->initialised) {

198 err = convolution_opencl_init(avctx);

199 if (err < 0)

200 goto fail;

201

202 if (!strcmp(avctx->filter->name, "convolution_opencl")) {

203 err = convolution_opencl_make_filter_params(avctx);

204 if (err < 0)

205 goto fail;

206 } else {

207 ctx->delta /= 255.0;

208 }

209

210 }

211

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

213 if (!output) {

214 err = AVERROR(ENOMEM);

215 goto fail;

216 }

217

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

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

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

221

222 if (!dst)

223 break;

224

225 if (!strcmp(avctx->filter->name, "convolution_opencl")) {

226 CL_SET_KERNEL_ARG(ctx->kernel, 0, cl_mem, &dst);

227 CL_SET_KERNEL_ARG(ctx->kernel, 1, cl_mem, &src);

228 CL_SET_KERNEL_ARG(ctx->kernel, 2, cl_int, &ctx->dims[p]);

229 CL_SET_KERNEL_ARG(ctx->kernel, 3, cl_mem, &ctx->matrix[p]);

230 CL_SET_KERNEL_ARG(ctx->kernel, 4, cl_float, &ctx->rdivs[p]);

231 CL_SET_KERNEL_ARG(ctx->kernel, 5, cl_float, &ctx->biases[p]);

232

233 err = ff_opencl_filter_work_size_from_image(avctx, global_work, output, p, 0);

234 if (err < 0)

235 goto fail;

236

237 av_log(avctx, AV_LOG_DEBUG, "Run kernel on plane %d "

238 "(%"SIZE_SPECIFIER "x%"SIZE_SPECIFIER ").\n",

239 p, global_work[0], global_work[1]);

240

241 cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,

242 global_work, NULL,

243 0, NULL, NULL);

244 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue "

245 "kernel: %d.\n", cle);

246 } else {

247 if (!(ctx->planes & (1 << p))) {

248 err = ff_opencl_filter_work_size_from_image(avctx, region, output, p, 0);

249 if (err < 0)

250 goto fail;

251

252 cle = clEnqueueCopyImage(ctx->command_queue, src, dst,

253 origin, origin, region, 0, NULL, NULL);

254 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to copy plane %d: %d.\n",

255 p, cle);

256 } else {

257 CL_SET_KERNEL_ARG(ctx->kernel, 0, cl_mem, &dst);

258 CL_SET_KERNEL_ARG(ctx->kernel, 1, cl_mem, &src);

259 CL_SET_KERNEL_ARG(ctx->kernel, 2, cl_float, &ctx->scale);

260 CL_SET_KERNEL_ARG(ctx->kernel, 3, cl_float, &ctx->delta);

261

262 err = ff_opencl_filter_work_size_from_image(avctx, global_work, output, p, 0);

263 if (err < 0)

264 goto fail;

265

266 av_log(avctx, AV_LOG_DEBUG, "Run kernel on plane %d "

267 "(%"SIZE_SPECIFIER "x%"SIZE_SPECIFIER ").\n",

268 p, global_work[0], global_work[1]);

269

270 cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,

271 global_work, NULL,

272 0, NULL, NULL);

273 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue "

274 "kernel: %d.\n", cle);

275 }

276 }

277 }

278

279 cle = clFinish(ctx->command_queue);

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

281

282 err = av_frame_copy_props(output, input);

283 if (err < 0)

284 goto fail;

285

286 av_frame_free(&input);

287

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

289 av_get_pix_fmt_name(output->format),

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

291

292 return ff_filter_frame(outlink, output);

293

294 fail:

295 clFinish(ctx->command_queue);

296 av_frame_free(&input);

297 av_frame_free(&output);

298 return err;

299 }

300

301 static av_cold void convolution_opencl_uninit(AVFilterContext *avctx)

302 {

303 ConvolutionOpenCLContext *ctx = avctx->priv;

304 cl_int cle;

305 int i;

306

307 for (i = 0; i < 4; i++) {

308 clReleaseMemObject(ctx->matrix[i]);

309 }

310

311 if (ctx->kernel) {

312 cle = clReleaseKernel(ctx->kernel);

313 if (cle != CL_SUCCESS)

314 av_log(avctx, AV_LOG_ERROR, "Failed to release "

315 "kernel: %d.\n", cle);

316 }

317

318 if (ctx->command_queue) {

319 cle = clReleaseCommandQueue(ctx->command_queue);

320 if (cle != CL_SUCCESS)

321 av_log(avctx, AV_LOG_ERROR, "Failed to release "

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

323 }

324

325 ff_opencl_filter_uninit(avctx);

326 }

327

328 static const AVFilterPad convolution_opencl_inputs[] = {

329 {

330 .name = "default",

331 .type = AVMEDIA_TYPE_VIDEO,

332 .filter_frame = &convolution_opencl_filter_frame,

333 .config_props = &ff_opencl_filter_config_input,

334 },

335 };

336

337 static const AVFilterPad convolution_opencl_outputs[] = {

338 {

339 .name = "default",

340 .type = AVMEDIA_TYPE_VIDEO,

341 .config_props = &ff_opencl_filter_config_output,

342 },

343 };

344

345 #define OFFSET(x) offsetof(ConvolutionOpenCLContext, x)

346 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)

347

348 #if CONFIG_CONVOLUTION_OPENCL_FILTER

349

350 static const AVOption convolution_opencl_options[] = {

351 { "0m", "set matrix for 2nd plane", OFFSET(matrix_str[0]), AV_OPT_TYPE_STRING, {.str="0 0 0 0 1 0 0 0 0"}, 0, 0, FLAGS },

352 { "1m", "set matrix for 2nd plane", OFFSET(matrix_str[1]), AV_OPT_TYPE_STRING, {.str="0 0 0 0 1 0 0 0 0"}, 0, 0, FLAGS },

353 { "2m", "set matrix for 3rd plane", OFFSET(matrix_str[2]), AV_OPT_TYPE_STRING, {.str="0 0 0 0 1 0 0 0 0"}, 0, 0, FLAGS },

354 { "3m", "set matrix for 4th plane", OFFSET(matrix_str[3]), AV_OPT_TYPE_STRING, {.str="0 0 0 0 1 0 0 0 0"}, 0, 0, FLAGS },

355 { "0rdiv", "set rdiv for 1nd plane", OFFSET(rdivs[0]), AV_OPT_TYPE_FLOAT, {.dbl=1.0}, 0.0, INT_MAX, FLAGS},

356 { "1rdiv", "set rdiv for 2nd plane", OFFSET(rdivs[1]), AV_OPT_TYPE_FLOAT, {.dbl=1.0}, 0.0, INT_MAX, FLAGS},

357 { "2rdiv", "set rdiv for 3rd plane", OFFSET(rdivs[2]), AV_OPT_TYPE_FLOAT, {.dbl=1.0}, 0.0, INT_MAX, FLAGS},

358 { "3rdiv", "set rdiv for 4th plane", OFFSET(rdivs[3]), AV_OPT_TYPE_FLOAT, {.dbl=1.0}, 0.0, INT_MAX, FLAGS},

359 { "0bias", "set bias for 1st plane", OFFSET(biases[0]), AV_OPT_TYPE_FLOAT, {.dbl=0.0}, 0.0, INT_MAX, FLAGS},

360 { "1bias", "set bias for 2nd plane", OFFSET(biases[1]), AV_OPT_TYPE_FLOAT, {.dbl=0.0}, 0.0, INT_MAX, FLAGS},

361 { "2bias", "set bias for 3rd plane", OFFSET(biases[2]), AV_OPT_TYPE_FLOAT, {.dbl=0.0}, 0.0, INT_MAX, FLAGS},

362 { "3bias", "set bias for 4th plane", OFFSET(biases[3]), AV_OPT_TYPE_FLOAT, {.dbl=0.0}, 0.0, INT_MAX, FLAGS},

363 { NULL }

364 };

365

366 AVFILTER_DEFINE_CLASS(convolution_opencl);

367

368 const AVFilter ff_vf_convolution_opencl = {

369 .name = "convolution_opencl",

370 .description = NULL_IF_CONFIG_SMALL("Apply convolution mask to input video"),

371 .priv_size = sizeof(ConvolutionOpenCLContext),

372 .priv_class = &convolution_opencl_class,

373 .init = &ff_opencl_filter_init,

374 .uninit = &convolution_opencl_uninit,

375 FILTER_INPUTS(convolution_opencl_inputs),

376 FILTER_OUTPUTS(convolution_opencl_outputs),

377 FILTER_SINGLE_PIXFMT(AV_PIX_FMT_OPENCL),

378 .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,

379 .flags = AVFILTER_FLAG_HWDEVICE,

380 };

381

382 #endif /* CONFIG_CONVOLUTION_OPENCL_FILTER */

383

384 #if CONFIG_SOBEL_OPENCL_FILTER

385

386 static const AVOption sobel_opencl_options[] = {

387 { "planes", "set planes to filter", OFFSET(planes), AV_OPT_TYPE_INT, {.i64=15}, 0, 15, FLAGS},

388 { "scale", "set scale", OFFSET(scale), AV_OPT_TYPE_FLOAT, {.dbl=1.0}, 0.0, 65535, FLAGS},

389 { "delta", "set delta", OFFSET(delta), AV_OPT_TYPE_FLOAT, {.dbl=0}, -65535, 65535, FLAGS},

390 { NULL }

391 };

392

393 AVFILTER_DEFINE_CLASS(sobel_opencl);

394

395 const AVFilter ff_vf_sobel_opencl = {

396 .name = "sobel_opencl",

397 .description = NULL_IF_CONFIG_SMALL("Apply sobel operator"),

398 .priv_size = sizeof(ConvolutionOpenCLContext),

399 .priv_class = &sobel_opencl_class,

400 .init = &ff_opencl_filter_init,

401 .uninit = &convolution_opencl_uninit,

402 FILTER_INPUTS(convolution_opencl_inputs),

403 FILTER_OUTPUTS(convolution_opencl_outputs),

404 FILTER_SINGLE_PIXFMT(AV_PIX_FMT_OPENCL),

405 .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,

406 .flags = AVFILTER_FLAG_HWDEVICE,

407 };

408

409 #endif /* CONFIG_SOBEL_OPENCL_FILTER */

410

411 #if CONFIG_PREWITT_OPENCL_FILTER

412

413 static const AVOption prewitt_opencl_options[] = {

414 { "planes", "set planes to filter", OFFSET(planes), AV_OPT_TYPE_INT, {.i64=15}, 0, 15, FLAGS},

415 { "scale", "set scale", OFFSET(scale), AV_OPT_TYPE_FLOAT, {.dbl=1.0}, 0.0, 65535, FLAGS},

416 { "delta", "set delta", OFFSET(delta), AV_OPT_TYPE_FLOAT, {.dbl=0}, -65535, 65535, FLAGS},

417 { NULL }

418 };

419

420 AVFILTER_DEFINE_CLASS(prewitt_opencl);

421

422 const AVFilter ff_vf_prewitt_opencl = {

423 .name = "prewitt_opencl",

424 .description = NULL_IF_CONFIG_SMALL("Apply prewitt operator"),

425 .priv_size = sizeof(ConvolutionOpenCLContext),

426 .priv_class = &prewitt_opencl_class,

427 .init = &ff_opencl_filter_init,

428 .uninit = &convolution_opencl_uninit,

429 FILTER_INPUTS(convolution_opencl_inputs),

430 FILTER_OUTPUTS(convolution_opencl_outputs),

431 FILTER_SINGLE_PIXFMT(AV_PIX_FMT_OPENCL),

432 .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,

433 .flags = AVFILTER_FLAG_HWDEVICE,

434 };

435

436 #endif /* CONFIG_PREWITT_OPENCL_FILTER */

437

438 #if CONFIG_ROBERTS_OPENCL_FILTER

439

440 static const AVOption roberts_opencl_options[] = {

441 { "planes", "set planes to filter", OFFSET(planes), AV_OPT_TYPE_INT, {.i64=15}, 0, 15, FLAGS},

442 { "scale", "set scale", OFFSET(scale), AV_OPT_TYPE_FLOAT, {.dbl=1.0}, 0.0, 65535, FLAGS},

443 { "delta", "set delta", OFFSET(delta), AV_OPT_TYPE_FLOAT, {.dbl=0}, -65535, 65535, FLAGS},

444 { NULL }

445 };

446

447 AVFILTER_DEFINE_CLASS(roberts_opencl);

448

449 const AVFilter ff_vf_roberts_opencl = {

450 .name = "roberts_opencl",

451 .description = NULL_IF_CONFIG_SMALL("Apply roberts operator"),

452 .priv_size = sizeof(ConvolutionOpenCLContext),

453 .priv_class = &roberts_opencl_class,

454 .init = &ff_opencl_filter_init,

455 .uninit = &convolution_opencl_uninit,

456 FILTER_INPUTS(convolution_opencl_inputs),

457 FILTER_OUTPUTS(convolution_opencl_outputs),

458 FILTER_SINGLE_PIXFMT(AV_PIX_FMT_OPENCL),

459 .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,

460 .flags = AVFILTER_FLAG_HWDEVICE,

461 };

462

463 #endif /* CONFIG_ROBERTS_OPENCL_FILTER */

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

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

CL_SET_KERNEL_ARG

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

set argument to specific Kernel.

Definition: opencl.h:61

convolution_opencl_filter_frame

static int convolution_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)

Definition: vf_convolution_opencl.c:177

ff_filter_frame

int ff_filter_frame(AVFilterLink *link, AVFrame *frame)

Send a frame of data to the next filter.

Definition: avfilter.c:1023

ConvolutionOpenCLContext::kernel

cl_kernel kernel

Definition: vf_convolution_opencl.c:42

matrix

Definition: vc1dsp.c:43

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

FILTER_INPUTS

#define FILTER_INPUTS(array)

Definition: filters.h:262

test::height

int height

Definition: vc1dsp.c:40

AVFrame

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

Definition: frame.h:389

pixdesc.h

opencl.h

AVOption

AVOption.

Definition: opt.h:429

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

ConvolutionOpenCLContext::ocf

OpenCLFilterContext ocf

Definition: vf_convolution_opencl.c:39

AVFilter::name

const char * name

Filter name.

Definition: avfilter.h:205

video.h

AVFilterLink

A link between two filters.

Definition: avfilter.h:557

av_malloc

#define av_malloc(s)

Definition: tableprint_vlc.h:30

ff_opencl_filter_work_size_from_image

int ff_opencl_filter_work_size_from_image(AVFilterContext *avctx, size_t *work_size, AVFrame *frame, int plane, int block_alignment)

Find the work size needed needed for a given plane of an image.

Definition: opencl.c:266

AVFilterContext::priv

void * priv

private data for use by the filter

Definition: avfilter.h:472

fail

#define fail()

Definition: checkasm.h:188

ConvolutionOpenCLContext::delta

cl_float delta

Definition: vf_convolution_opencl.c:55

ConvolutionOpenCLContext::dims

cl_int dims[4]

Definition: vf_convolution_opencl.c:49

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

convolution_opencl_inputs

static const AVFilterPad convolution_opencl_inputs[]

Definition: vf_convolution_opencl.c:328

AVFilterPad

A filter pad used for either input or output.

Definition: filters.h:38

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

convolution_opencl_outputs

static const AVFilterPad convolution_opencl_outputs[]

Definition: vf_convolution_opencl.c:337

ConvolutionOpenCLContext::matrix_sizes

cl_int matrix_sizes[4]

Definition: vf_convolution_opencl.c:48

float

Definition: af_crystalizer.c:122

av_strtok

char * av_strtok(char *s, const char *delim, char **saveptr)

Split the string into several tokens which can be accessed by successive calls to av_strtok().

Definition: avstring.c:178

av_assert0

#define av_assert0(cond)

assert() equivalent, that is always enabled.

Definition: avassert.h:40

filters.h

AV_LOG_DEBUG

#define AV_LOG_DEBUG

Stuff which is only useful for libav* developers.

Definition: log.h:201

ConvolutionOpenCLContext

Definition: vf_convolution_opencl.c:38

ctx

AVFormatContext * ctx

Definition: movenc.c:49

FILTER_OUTPUTS

#define FILTER_OUTPUTS(array)

Definition: filters.h:263

arg

const char * arg

Definition: jacosubdec.c:67

NULL

#define NULL

Definition: coverity.c:32

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

ConvolutionOpenCLContext::matrix_str

char * matrix_str[4]

Definition: vf_convolution_opencl.c:45

ConvolutionOpenCLContext::command_queue

cl_command_queue command_queue

Definition: vf_convolution_opencl.c:43

AV_PIX_FMT_OPENCL

@ AV_PIX_FMT_OPENCL

Hardware surfaces for OpenCL.

Definition: pixfmt.h:358

ConvolutionOpenCLContext::rdivs

cl_float rdivs[4]

Definition: vf_convolution_opencl.c:50

AVFILTER_DEFINE_CLASS

#define AVFILTER_DEFINE_CLASS(fname)

Definition: filters.h:273

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

test::width

int width

Definition: vc1dsp.c:39

init

int(* init)(AVBSFContext *ctx)

Definition: dts2pts.c:368

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

dst

uint8_t ptrdiff_t const uint8_t ptrdiff_t int intptr_t intptr_t int int16_t * dst

Definition: dsp.h:83

ff_source_convolution_cl

const char * ff_source_convolution_cl

ConvolutionOpenCLContext::initialised

int initialised

Definition: vf_convolution_opencl.c:41

AVFILTER_FLAG_HWDEVICE

#define AVFILTER_FLAG_HWDEVICE

The filter can create hardware frames using AVFilterContext.hw_device_ctx.

Definition: avfilter.h:173

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

AV_OPT_TYPE_FLOAT

@ AV_OPT_TYPE_FLOAT

Underlying C type is float.

Definition: opt.h:271

uninit

static void uninit(AVBSFContext *ctx)

Definition: pcm_rechunk.c:68

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

Definition: cbs_h2645.c:256

AVFilterLink::w

int w

agreed upon image width

Definition: avfilter.h:569

common.h

delta

float delta

Definition: vorbis_enc_data.h:430

AVFilterPad::name

const char * name

Pad name.

Definition: filters.h:44

AVFilter

Filter definition.

Definition: avfilter.h:201

ConvolutionOpenCLContext::planes

cl_int planes

Definition: vf_convolution_opencl.c:53

ff_opencl_filter_init

int ff_opencl_filter_init(AVFilterContext *avctx)

Initialise an OpenCL filter context.

Definition: opencl.c:135

ff_vf_roberts_opencl

const AVFilter ff_vf_roberts_opencl

ConvolutionOpenCLContext::scale

cl_float scale

Definition: vf_convolution_opencl.c:54

SIZE_SPECIFIER

#define SIZE_SPECIFIER

Definition: internal.h:129

ff_vf_convolution_opencl

const AVFilter ff_vf_convolution_opencl

ConvolutionOpenCLContext::biases

cl_float biases[4]

Definition: vf_convolution_opencl.c:51

OFFSET

#define OFFSET(x)

Definition: vf_convolution_opencl.c:345

planes

static const struct @455 planes[]

AVFilterLink::h

int h

agreed upon image height

Definition: avfilter.h:570

buffer

the frame and frame reference mechanism is intended to as much as expensive copies of that data while still allowing the filters to produce correct results The data is stored in buffers represented by AVFrame structures Several references can point to the same frame buffer

Definition: filter_design.txt:49

ff_vf_sobel_opencl

const AVFilter ff_vf_sobel_opencl

AV_OPT_TYPE_INT

@ AV_OPT_TYPE_INT

Underlying C type is int.

Definition: opt.h:259

avfilter.h

convolution_opencl_init

static int convolution_opencl_init(AVFilterContext *avctx)

Definition: vf_convolution_opencl.c:59

OpenCLFilterContext