1 /*
2 * Copyright (c) 2018 Dylan Fernando
3 * Copyright (c) 2018 Danil Iashchenko
4 *
5 * This file is part of FFmpeg.
6 *
7 * FFmpeg is free software; you can redistribute it and/or
8 * modify it under the terms of the GNU Lesser General Public
9 * License as published by the Free Software Foundation; either
10 * version 2.1 of the License, or (at your option) any later version.
11 *
12 * FFmpeg is distributed in the hope that it will be useful,
13 * but WITHOUT ANY WARRANTY; without even the implied warranty of
14 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
15 * Lesser General Public License for more details.
16 *
17 * You should have received a copy of the GNU Lesser General Public
18 * License along with FFmpeg; if not, write to the Free Software
19 * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
20 */
21
22 #include "config_components.h"
23
27
34
37
42
46
52
54
55
57 {
59 cl_int cle;
60 int err;
61
63 if (err < 0)
65
66 ctx->command_queue = clCreateCommandQueue(
ctx->ocf.hwctx->context,
67 ctx->ocf.hwctx->device_id,
68 0, &cle);
70 "command queue %d.\n", cle);
71
72 ctx->kernel_horiz = clCreateKernel(
ctx->ocf.program,
"avgblur_horiz", &cle);
74 "kernel %d.\n", cle);
75
76 ctx->kernel_vert = clCreateKernel(
ctx->ocf.program,
"avgblur_vert", &cle);
78 "kernel %d.\n", cle);
79
81 return 0;
82
84 if (
ctx->command_queue)
85 clReleaseCommandQueue(
ctx->command_queue);
86 if (
ctx->kernel_horiz)
87 clReleaseKernel(
ctx->kernel_horiz);
89 clReleaseKernel(
ctx->kernel_vert);
90 return err;
91 }
92
93
95 {
99
100 if (
s->radiusV <= 0) {
101 s->radiusV =
s->radiusH;
102 }
103
106 }
107 return 0;
108 }
109
110
112 {
116
121
122 if (err != 0) {
124 "filter params: %d.\n", err);
125 return err;
126 }
127
128 s->radius[
Y] =
s->luma_param.radius;
129 s->radius[
U] =
s->radius[
V] =
s->chroma_param.radius;
130 s->radius[
A] =
s->alpha_param.radius;
131
132 s->power[
Y] =
s->luma_param.power;
133 s->power[
U] =
s->power[
V] =
s->chroma_param.power;
134 s->power[
A] =
s->alpha_param.power;
135
137 if (
s->power[
i] == 0) {
140 }
141 }
142
143 return 0;
144 }
145
146
148 {
154 cl_int cle;
155 size_t global_work[2];
157 int err,
p, radius_x, radius_y,
i;
158
162
163 if (!
input->hw_frames_ctx)
165
166 if (!
ctx->initialised) {
168 if (err < 0)
170
171 if (!strcmp(avctx->
filter->
name,
"avgblur_opencl")) {
173 if (err < 0)
175 }
else if (!strcmp(avctx->
filter->
name,
"boxblur_opencl")) {
177 if (err < 0)
179 }
180
181 }
182
187 }
189 if (!intermediate) {
192 }
193
197 inter = (cl_mem)intermediate->
data[
p];
198
200 break;
201
202 radius_x =
ctx->radiusH;
203 radius_y =
ctx->radiusV;
204
205 if (!(
ctx->planes & (1 <<
p))) {
206 radius_x = 0;
207 radius_y = 0;
208 }
209
210 for (
i = 0;
i <
ctx->power[
p];
i++) {
213 if (!strcmp(avctx->
filter->
name,
"avgblur_opencl")) {
215 }
else if (!strcmp(avctx->
filter->
name,
"boxblur_opencl")) {
217 }
218
220 i == 0 ? intermediate :
output,
p, 0);
221 if (err < 0)
223
226 p, global_work[0], global_work[1]);
227
228 cle = clEnqueueNDRangeKernel(
ctx->command_queue,
ctx->kernel_horiz, 2,
NULL,
232 "kernel: %d.\n", cle);
233
235 i == 0 ?
output : intermediate,
p, 0);
236
239
240 if (!strcmp(avctx->
filter->
name,
"avgblur_opencl")) {
242 }
else if (!strcmp(avctx->
filter->
name,
"boxblur_opencl")) {
244 }
245
246 cle = clEnqueueNDRangeKernel(
ctx->command_queue,
ctx->kernel_vert, 2,
NULL,
250 "kernel: %d.\n", cle);
251 }
252 }
253
254 cle = clFinish(
ctx->command_queue);
256
258 if (err < 0)
260
263
267
269
271 clFinish(
ctx->command_queue);
275 return err;
276 }
277
278
280 {
282 cl_int cle;
283
284 if (
ctx->kernel_horiz) {
285 cle = clReleaseKernel(
ctx->kernel_horiz);
286 if (cle != CL_SUCCESS)
288 "kernel: %d.\n", cle);
289 }
290
291 if (
ctx->kernel_vert) {
292 cle = clReleaseKernel(
ctx->kernel_vert);
293 if (cle != CL_SUCCESS)
295 "kernel: %d.\n", cle);
296 }
297
298 if (
ctx->command_queue) {
299 cle = clReleaseCommandQueue(
ctx->command_queue);
300 if (cle != CL_SUCCESS)
302 "command queue: %d.\n", cle);
303 }
304
306 }
307
308
310 {
315 },
316 };
317
318
320 {
324 },
325 };
326
327
328 #define OFFSET(x) offsetof(AverageBlurOpenCLContext, x)
329 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
330
331 #if CONFIG_AVGBLUR_OPENCL_FILTER
332
333 static const AVOption avgblur_opencl_options[] = {
338 };
339
341
342
344 .
p.
name =
"avgblur_opencl",
346 .p.priv_class = &avgblur_opencl_class,
354 };
355
356 #endif /* CONFIG_AVGBLUR_OPENCL_FILTER */
357
358
359 #if CONFIG_BOXBLUR_OPENCL_FILTER
360
361 static const AVOption boxblur_opencl_options[] = {
364 {
"luma_power",
"How many times should the boxblur be applied to luma",
OFFSET(luma_param.power),
AV_OPT_TYPE_INT, {.i64=2}, 0, INT_MAX, .flags =
FLAGS },
365 {
"lp",
"How many times should the boxblur be applied to luma",
OFFSET(luma_param.power),
AV_OPT_TYPE_INT, {.i64=2}, 0, INT_MAX, .flags =
FLAGS },
366
369 {
"chroma_power",
"How many times should the boxblur be applied to chroma",
OFFSET(chroma_param.power),
AV_OPT_TYPE_INT, {.i64=-1}, -1, INT_MAX, .flags =
FLAGS },
370 {
"cp",
"How many times should the boxblur be applied to chroma",
OFFSET(chroma_param.power),
AV_OPT_TYPE_INT, {.i64=-1}, -1, INT_MAX, .flags =
FLAGS },
371
374 {
"alpha_power",
"How many times should the boxblur be applied to alpha",
OFFSET(alpha_param.power),
AV_OPT_TYPE_INT, {.i64=-1}, -1, INT_MAX, .flags =
FLAGS },
375 {
"ap",
"How many times should the boxblur be applied to alpha",
OFFSET(alpha_param.power),
AV_OPT_TYPE_INT, {.i64=-1}, -1, INT_MAX, .flags =
FLAGS },
376
378 };
379
381
383 .
p.
name =
"boxblur_opencl",
385 .p.priv_class = &boxblur_opencl_class,
394 };
395
396 #endif /* CONFIG_BOXBLUR_OPENCL_FILTER */