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/common.h"

24 #include "libavutil/imgutils.h"

25 #include "libavutil/mem.h"

26 #include "libavutil/opt.h"

27 #include "libavutil/pixdesc.h"

28 #include "libavutil/avstring.h"

31 #include "avfilter.h"

32 #include "internal.h"

33 #include "opencl.h"

34 #include "opencl_source.h"

35 #include "video.h"

37 typedef struct ConvolutionOpenCLContext {

38 OpenCLFilterContext ocf;

40 int initialised;

41 cl_kernel kernel;

42 cl_command_queue command_queue;

44 char *matrix_str[4];

46 cl_mem matrix[4];

47 cl_int matrix_sizes[4];

48 cl_int dims[4];

49 cl_float rdivs[4];

50 cl_float biases[4];

52 cl_int planes;

53 cl_float scale;

54 cl_float delta;

56 } ConvolutionOpenCLContext;

58 static int convolution_opencl_init(AVFilterContext *avctx)

59 {

60 ConvolutionOpenCLContext *ctx = avctx->priv;

61 const char *kernel_name;

62 cl_int cle;

63 int err;

65 err = ff_opencl_filter_load_program(avctx, &ff_opencl_source_convolution, 1);

66 if (err < 0)

67 goto fail;

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

70 ctx->ocf.hwctx->device_id,

71 0, &cle);

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

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

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

76 kernel_name = "convolution_global";

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

78 kernel_name = "sobel_global";

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

80 kernel_name = "prewitt_global";

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

82 kernel_name = "roberts_global";

83 }

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

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

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

88 ctx->initialised = 1;

89 return 0;

91 fail:

92 if (ctx->command_queue)

93 clReleaseCommandQueue(ctx->command_queue);

94 if (ctx->kernel)

95 clReleaseKernel(ctx->kernel);

96 return err;

97 }

100

101 static int convolution_opencl_make_filter_params(AVFilterContext *avctx)

102 {

103 ConvolutionOpenCLContext *ctx = avctx->priv;

104 float *matrix = NULL;

105 size_t matrix_bytes;

106 cl_mem buffer;

107 cl_int cle;

108 int i, j;

109 int sscanf_err;

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

111 float input_matrix[4][49];

112

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

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

115 }

116

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

118 p = ctx->matrix_str[i];

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

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

121 if (!arg) {

122 break;

123 }

124 p = NULL;

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

126 if (sscanf_err != 1) {

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

128 return AVERROR(EINVAL);

129 }

130 ctx->matrix_sizes[i]++;

131 }

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

133 ctx->dims[i] = 3;

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

135 ctx->dims[i] = 5;

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

137 ctx->dims[i] = 7;

138 } else {

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

140 return AVERROR(EINVAL);

141 }

142

143 }

144

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

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

147 matrix = av_malloc(matrix_bytes);

148 if (!matrix) {

149 av_freep(&matrix);

150 return AVERROR(ENOMEM);

151 }

152

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

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

155

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

157 CL_MEM_READ_ONLY |

158 CL_MEM_COPY_HOST_PTR |

159 CL_MEM_HOST_NO_ACCESS,

160 matrix_bytes, matrix, &cle);

161 if (!buffer) {

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

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

164 av_freep(&matrix);

165 return AVERROR(EIO);

166 }

167 ctx->matrix[j] = buffer;

168 av_freep(&matrix);

169 }

170

171 return 0;

172 }

173

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

175 {

176 AVFilterContext *avctx = inlink->dst;

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

178 ConvolutionOpenCLContext *ctx = avctx->priv;

179 AVFrame *output = NULL;

180 cl_int cle;

181 size_t global_work[2];

182 cl_mem src, dst;

183 int err, p;

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

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

186

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

188 av_get_pix_fmt_name(input->format),

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

190

191 if (!input->hw_frames_ctx)

192 return AVERROR(EINVAL);

193

194 if (!ctx->initialised) {

195 err = convolution_opencl_init(avctx);

196 if (err < 0)

197 goto fail;

198

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

200 err = convolution_opencl_make_filter_params(avctx);

201 if (err < 0)

202 goto fail;

203 } else {

204 ctx->delta /= 255.0;

205 }

206

207 }

208

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

210 if (!output) {

211 err = AVERROR(ENOMEM);

212 goto fail;

213 }

214

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

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

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

218

219 if (!dst)

220 break;

221

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

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

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

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

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

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

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

229

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

231 if (err < 0)

232 goto fail;

233

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

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

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

237

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

239 global_work, NULL,

240 0, NULL, NULL);

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

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

243 } else {

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

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

246 if (err < 0)

247 goto fail;

248

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

250 origin, origin, region, 0, NULL, NULL);

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

252 p, cle);

253 } else {

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

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

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

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

258

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

260 if (err < 0)

261 goto fail;

262

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

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

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

266

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

268 global_work, NULL,

269 0, NULL, NULL);

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

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

272 }

273 }

274 }

275

276 cle = clFinish(ctx->command_queue);

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

278

279 err = av_frame_copy_props(output, input);

280 if (err < 0)

281 goto fail;

282

283 av_frame_free(&input);

284

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

286 av_get_pix_fmt_name(output->format),

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

288

289 return ff_filter_frame(outlink, output);

290

291 fail:

292 clFinish(ctx->command_queue);

293 av_frame_free(&input);

294 av_frame_free(&output);

295 return err;

296 }

297

298 static av_cold void convolution_opencl_uninit(AVFilterContext *avctx)

299 {

300 ConvolutionOpenCLContext *ctx = avctx->priv;

301 cl_int cle;

302 int i;

303

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

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

306 }

307

308 if (ctx->kernel) {

309 cle = clReleaseKernel(ctx->kernel);

310 if (cle != CL_SUCCESS)

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

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

313 }

314

315 if (ctx->command_queue) {

316 cle = clReleaseCommandQueue(ctx->command_queue);

317 if (cle != CL_SUCCESS)

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

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

320 }

321

322 ff_opencl_filter_uninit(avctx);

323 }

324

325 static const AVFilterPad convolution_opencl_inputs[] = {

326 {

327 .name = "default",

328 .type = AVMEDIA_TYPE_VIDEO,

329 .filter_frame = &convolution_opencl_filter_frame,

330 .config_props = &ff_opencl_filter_config_input,

331 },

332 };

333

334 static const AVFilterPad convolution_opencl_outputs[] = {

335 {

336 .name = "default",

337 .type = AVMEDIA_TYPE_VIDEO,

338 .config_props = &ff_opencl_filter_config_output,

339 },

340 };

341

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

343 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)

344

345 #if CONFIG_CONVOLUTION_OPENCL_FILTER

346

347 static const AVOption convolution_opencl_options[] = {

348 { "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 },

349 { "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 },

350 { "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 },

351 { "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 },

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

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

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

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

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

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

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

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

360 { NULL }

361 };

362

363 AVFILTER_DEFINE_CLASS(convolution_opencl);

364

365 const AVFilter ff_vf_convolution_opencl = {

366 .name = "convolution_opencl",

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

368 .priv_size = sizeof(ConvolutionOpenCLContext),

369 .priv_class = &convolution_opencl_class,

370 .init = &ff_opencl_filter_init,

371 .uninit = &convolution_opencl_uninit,

372 FILTER_INPUTS(convolution_opencl_inputs),

373 FILTER_OUTPUTS(convolution_opencl_outputs),

374 FILTER_SINGLE_PIXFMT(AV_PIX_FMT_OPENCL),

375 .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,

376 };

377

378 #endif /* CONFIG_CONVOLUTION_OPENCL_FILTER */

379

380 #if CONFIG_SOBEL_OPENCL_FILTER

381

382 static const AVOption sobel_opencl_options[] = {

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

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

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

386 { NULL }

387 };

388

389 AVFILTER_DEFINE_CLASS(sobel_opencl);

390

391 const AVFilter ff_vf_sobel_opencl = {

392 .name = "sobel_opencl",

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

394 .priv_size = sizeof(ConvolutionOpenCLContext),

395 .priv_class = &sobel_opencl_class,

396 .init = &ff_opencl_filter_init,

397 .uninit = &convolution_opencl_uninit,

398 FILTER_INPUTS(convolution_opencl_inputs),

399 FILTER_OUTPUTS(convolution_opencl_outputs),

400 FILTER_SINGLE_PIXFMT(AV_PIX_FMT_OPENCL),

401 .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,

402 };

403

404 #endif /* CONFIG_SOBEL_OPENCL_FILTER */

405

406 #if CONFIG_PREWITT_OPENCL_FILTER

407

408 static const AVOption prewitt_opencl_options[] = {

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

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

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

412 { NULL }

413 };

414

415 AVFILTER_DEFINE_CLASS(prewitt_opencl);

416

417 const AVFilter ff_vf_prewitt_opencl = {

418 .name = "prewitt_opencl",

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

420 .priv_size = sizeof(ConvolutionOpenCLContext),

421 .priv_class = &prewitt_opencl_class,

422 .init = &ff_opencl_filter_init,

423 .uninit = &convolution_opencl_uninit,

424 FILTER_INPUTS(convolution_opencl_inputs),

425 FILTER_OUTPUTS(convolution_opencl_outputs),

426 FILTER_SINGLE_PIXFMT(AV_PIX_FMT_OPENCL),

427 .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,

428 };

429

430 #endif /* CONFIG_PREWITT_OPENCL_FILTER */

431

432 #if CONFIG_ROBERTS_OPENCL_FILTER

433

434 static const AVOption roberts_opencl_options[] = {

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

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

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

438 { NULL }

439 };

440

441 AVFILTER_DEFINE_CLASS(roberts_opencl);

442

443 const AVFilter ff_vf_roberts_opencl = {

444 .name = "roberts_opencl",

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

446 .priv_size = sizeof(ConvolutionOpenCLContext),

447 .priv_class = &roberts_opencl_class,

448 .init = &ff_opencl_filter_init,

449 .uninit = &convolution_opencl_uninit,

450 FILTER_INPUTS(convolution_opencl_inputs),

451 FILTER_OUTPUTS(convolution_opencl_outputs),

452 FILTER_SINGLE_PIXFMT(AV_PIX_FMT_OPENCL),

453 .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,

454 };

455

456 #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:101

planes

static const struct @346 planes[]

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

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

ConvolutionOpenCLContext::kernel

cl_kernel kernel

Definition: vf_convolution_opencl.c:41

matrix

Definition: vc1dsp.c:42

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

AVFrame

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

Definition: frame.h:330

pixdesc.h

opencl.h

AVOption

AVOption.

Definition: opt.h:251

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

ConvolutionOpenCLContext::ocf

OpenCLFilterContext ocf

Definition: vf_convolution_opencl.c:38

AVFilter::name

const char * name

Filter name.

Definition: avfilter.h:165

video.h

AVFilterLink

A link between two filters.

Definition: avfilter.h:522

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

AVFilterContext::priv

void * priv

private data for use by the filter

Definition: avfilter.h:407

fail

#define fail()

Definition: checkasm.h:134

ConvolutionOpenCLContext::delta

cl_float delta

Definition: vf_convolution_opencl.c:54

scale

static av_always_inline float scale(float x, float s)

Definition: vf_v360.c:1389

ff_opencl_source_convolution

const char * ff_opencl_source_convolution

ConvolutionOpenCLContext::dims

cl_int dims[4]

Definition: vf_convolution_opencl.c:48

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

convolution_opencl_inputs

static const AVFilterPad convolution_opencl_inputs[]

Definition: vf_convolution_opencl.c:325

AVFilterPad

A filter pad used for either input or output.

Definition: internal.h:49

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

ConvolutionOpenCLContext::matrix_sizes

cl_int matrix_sizes[4]

Definition: vf_convolution_opencl.c:47

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

init

int(* init)(AVBSFContext *ctx)

Definition: dts2pts_bsf.c:365

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

ctx

AVFormatContext * ctx

Definition: movenc.c:48

FILTER_INPUTS

#define FILTER_INPUTS(array)

Definition: internal.h:194

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

ConvolutionOpenCLContext::matrix_str

char * matrix_str[4]

Definition: vf_convolution_opencl.c:44

ConvolutionOpenCLContext::command_queue

cl_command_queue command_queue

Definition: vf_convolution_opencl.c:42

AV_PIX_FMT_OPENCL

@ AV_PIX_FMT_OPENCL

Hardware surfaces for OpenCL.

Definition: pixfmt.h:355

ConvolutionOpenCLContext::rdivs

cl_float rdivs[4]

Definition: vf_convolution_opencl.c:49

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

ConvolutionOpenCLContext::initialised

int initialised

Definition: vf_convolution_opencl.c:40

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

AVFILTER_DEFINE_CLASS

#define AVFILTER_DEFINE_CLASS(fname)

Definition: internal.h:329

AV_OPT_TYPE_FLOAT

@ AV_OPT_TYPE_FLOAT

Definition: opt.h:228

FILTER_SINGLE_PIXFMT

#define FILTER_SINGLE_PIXFMT(pix_fmt_)

Definition: internal.h:184

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

Definition: cbs_h2645.c:269

AVFilterLink::w

int w

agreed upon image width

Definition: avfilter.h:532

common.h

delta

float delta

Definition: vorbis_enc_data.h:430

AVFilterPad::name

const char * name

Pad name.

Definition: internal.h:55

AVFilter

Filter definition.

Definition: avfilter.h:161

ConvolutionOpenCLContext::planes

cl_int planes

Definition: vf_convolution_opencl.c:52

ff_opencl_filter_init

int ff_opencl_filter_init(AVFilterContext *avctx)

Initialise an OpenCL filter context.

Definition: opencl.c:133

ff_vf_roberts_opencl

const AVFilter ff_vf_roberts_opencl

ConvolutionOpenCLContext::scale

cl_float scale

Definition: vf_convolution_opencl.c:53

SIZE_SPECIFIER

#define SIZE_SPECIFIER

Definition: internal.h:150

ff_vf_convolution_opencl

const AVFilter ff_vf_convolution_opencl

ConvolutionOpenCLContext::biases

cl_float biases[4]

Definition: vf_convolution_opencl.c:50

OFFSET

#define OFFSET(x)

Definition: vf_convolution_opencl.c:342

AVFilterLink::h

int h

agreed upon image height

Definition: avfilter.h:533

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

Definition: opt.h:225

avfilter.h

convolution_opencl_init

static int convolution_opencl_init(AVFilterContext *avctx)

Definition: vf_convolution_opencl.c:58

OpenCLFilterContext

Definition: opencl.h:36