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
22
29
32
36
38
43
47
51 {
53 cl_int cle;
55 const char *kernel;
57 int err,
i, main_planes, overlay_planes;
58
61
62 main_planes = overlay_planes = 0;
64 main_planes =
FFMAX(main_planes,
67 overlay_planes =
FFMAX(overlay_planes,
69
70 ctx->nb_planes = main_planes;
73
74 if (
ctx->x_position %
ctx->x_subsample ||
75 ctx->y_position %
ctx->y_subsample) {
77 "does not match subsampling (%d, %d).\n",
78 ctx->x_position,
ctx->y_position,
79 ctx->x_subsample,
ctx->y_subsample);
80 }
81
82 if (main_planes == overlay_planes) {
84 kernel = "overlay_no_alpha";
85 else
86 kernel = "overlay_internal_alpha";
87 ctx->alpha_separate = 0;
88 } else {
89 kernel = "overlay_external_alpha";
90 ctx->alpha_separate = 1;
91 }
92
94
96 if (err < 0)
98
99 ctx->command_queue = clCreateCommandQueue(
ctx->ocf.hwctx->context,
100 ctx->ocf.hwctx->device_id,
101 0, &cle);
103 "command queue %d.\n", cle);
104
105 ctx->kernel = clCreateKernel(
ctx->ocf.program, kernel, &cle);
107
108 ctx->initialised = 1;
109 return 0;
110
112 if (
ctx->command_queue)
113 clReleaseCommandQueue(
ctx->command_queue);
115 clReleaseKernel(
ctx->kernel);
116 return err;
117 }
118
120 {
124 AVFrame *input_main, *input_overlay;
126 cl_mem mem;
127 cl_int cle, x, y;
128 size_t global_work[2];
129 int kernel_arg = 0;
130 int err, plane;
131
133 if (err < 0)
134 return err;
136 if (err < 0)
137 return err;
138
139 if (!
ctx->initialised) {
144
146 overlay_fc->sw_format);
147 if (err < 0)
148 return err;
149 }
150
155 }
156
157 for (plane = 0; plane <
ctx->nb_planes; plane++) {
158 kernel_arg = 0;
159
160 mem = (cl_mem)
output->data[plane];
162 kernel_arg++;
163
164 mem = (cl_mem)input_main->
data[plane];
166 kernel_arg++;
167
168 mem = (cl_mem)input_overlay->
data[plane];
170 kernel_arg++;
171
172 if (
ctx->alpha_separate) {
173 mem = (cl_mem)input_overlay->
data[
ctx->nb_planes];
175 kernel_arg++;
176 }
177
178 x =
ctx->x_position / (plane == 0 ? 1 :
ctx->x_subsample);
179 y =
ctx->y_position / (plane == 0 ? 1 :
ctx->y_subsample);
180
182 kernel_arg++;
184 kernel_arg++;
185
186 if (
ctx->alpha_separate) {
187 cl_int alpha_adj_x = plane == 0 ? 1 :
ctx->x_subsample;
188 cl_int alpha_adj_y = plane == 0 ? 1 :
ctx->y_subsample;
189
191 kernel_arg++;
193 kernel_arg++;
194 }
195
198 if (err < 0)
200
201 cle = clEnqueueNDRangeKernel(
ctx->command_queue,
ctx->kernel, 2,
NULL,
204 "for plane %d: %d.\n", plane, cle);
205 }
206
207 cle = clFinish(
ctx->command_queue);
209
211
215
217
220 return err;
221 }
222
224 {
227 int err;
228
230 if (err < 0)
231 return err;
232
234 if (err < 0)
235 return err;
236
238 }
239
241 {
243
245
247 }
248
250 {
252
254 }
255
257 {
259 cl_int cle;
260
262 cle = clReleaseKernel(
ctx->kernel);
263 if (cle != CL_SUCCESS)
265 "kernel: %d.\n", cle);
266 }
267
268 if (
ctx->command_queue) {
269 cle = clReleaseCommandQueue(
ctx->command_queue);
270 if (cle != CL_SUCCESS)
272 "command queue: %d.\n", cle);
273 }
274
276
278 }
279
280 #define OFFSET(x) offsetof(OverlayOpenCLContext, x)
281 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
283 { "x", "Overlay x position",
285 { "y", "Overlay y position",
288 };
289
291
293 {
297 },
298 {
299 .name = "overlay",
302 },
303 };
304
306 {
310 },
311 };
312
314 .
name =
"overlay_opencl",
317 .priv_class = &overlay_opencl_class,
325 };