avfilter: remove av_opt_set_defaults() from process_options()
[ffmpeg.git] / libavfilter / unsharp_opencl.c
1 /*
2  * Copyright (C) 2013 Wei Gao <weigao@multicorewareinc.com>
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  */
20
21 /**
22  * @file
23  * unsharp input video
24  */
25
26 #include "unsharp_opencl.h"
27 #include "libavutil/common.h"
28 #include "libavutil/opencl_internal.h"
29
30 #define PLANE_NUM 3
31
32 static inline void add_mask_counter(uint32_t *dst, uint32_t *counter1, uint32_t *counter2, int len)
33 {
34     int i;
35     for (i = 0; i < len; i++) {
36         dst[i] = counter1[i] + counter2[i];
37     }
38 }
39
40 static int compute_mask(int step, uint32_t *mask)
41 {
42     int i, z, ret = 0;
43     int counter_size = sizeof(uint32_t) * (2 * step + 1);
44     uint32_t *temp1_counter, *temp2_counter, **counter;
45     temp1_counter = av_mallocz(counter_size);
46     if (!temp1_counter) {
47         ret = AVERROR(ENOMEM);
48         goto end;
49     }
50     temp2_counter = av_mallocz(counter_size);
51     if (!temp2_counter) {
52         ret = AVERROR(ENOMEM);
53         goto end;
54     }
55     counter = av_mallocz(sizeof(uint32_t *) * (2 * step + 1));
56     if (!counter) {
57         ret = AVERROR(ENOMEM);
58         goto end;
59     }
60     for (i = 0; i < 2 * step + 1; i++) {
61         counter[i] = av_mallocz(counter_size);
62         if (!counter[i]) {
63             ret = AVERROR(ENOMEM);
64             goto end;
65         }
66     }
67     for (i = 0; i < 2 * step + 1; i++) {
68         memset(temp1_counter, 0, counter_size);
69         temp1_counter[i] = 1;
70         for (z = 0; z < step * 2; z += 2) {
71             add_mask_counter(temp2_counter, counter[z], temp1_counter, step * 2);
72             memcpy(counter[z], temp1_counter, counter_size);
73             add_mask_counter(temp1_counter, counter[z + 1], temp2_counter, step * 2);
74             memcpy(counter[z + 1], temp2_counter, counter_size);
75         }
76     }
77     memcpy(mask, temp1_counter, counter_size);
78 end:
79     av_freep(&temp1_counter);
80     av_freep(&temp2_counter);
81     for (i = 0; i < 2 * step + 1; i++) {
82         av_freep(&counter[i]);
83     }
84     av_freep(&counter);
85     return ret;
86 }
87
88 static int compute_mask_matrix(cl_mem cl_mask_matrix, int step_x, int step_y)
89 {
90     int i, j, ret = 0;
91     uint32_t *mask_matrix, *mask_x, *mask_y;
92     size_t size_matrix = sizeof(uint32_t) * (2 * step_x + 1) * (2 * step_y + 1);
93     mask_x = av_mallocz(sizeof(uint32_t) * (2 * step_x + 1));
94     if (!mask_x) {
95         ret = AVERROR(ENOMEM);
96         goto end;
97     }
98     mask_y = av_mallocz(sizeof(uint32_t) * (2 * step_y + 1));
99     if (!mask_y) {
100         ret = AVERROR(ENOMEM);
101         goto end;
102     }
103     mask_matrix = av_mallocz(size_matrix);
104     if (!mask_matrix) {
105         ret = AVERROR(ENOMEM);
106         goto end;
107     }
108     ret = compute_mask(step_x, mask_x);
109     if (ret < 0)
110         goto end;
111     ret = compute_mask(step_y, mask_y);
112     if (ret < 0)
113         goto end;
114     for (j = 0; j < 2 * step_y + 1; j++) {
115         for (i = 0; i < 2 * step_x + 1; i++) {
116             mask_matrix[i + j * (2 * step_x + 1)] = mask_y[j] * mask_x[i];
117         }
118     }
119     ret = av_opencl_buffer_write(cl_mask_matrix, (uint8_t *)mask_matrix, size_matrix);
120 end:
121     av_freep(&mask_x);
122     av_freep(&mask_y);
123     av_freep(&mask_matrix);
124     return ret;
125 }
126
127 static int generate_mask(AVFilterContext *ctx)
128 {
129     UnsharpContext *unsharp = ctx->priv;
130     int i, ret = 0, step_x[2], step_y[2];
131     cl_mem mask_matrix[2];
132     mask_matrix[0] = unsharp->opencl_ctx.cl_luma_mask;
133     mask_matrix[1] = unsharp->opencl_ctx.cl_chroma_mask;
134     step_x[0] = unsharp->luma.steps_x;
135     step_x[1] = unsharp->chroma.steps_x;
136     step_y[0] = unsharp->luma.steps_y;
137     step_y[1] = unsharp->chroma.steps_y;
138     if (!mask_matrix[0] || !mask_matrix[1]) {
139         av_log(ctx, AV_LOG_ERROR, "Luma mask and chroma mask should not be NULL\n");
140         return AVERROR(EINVAL);
141     }
142     for (i = 0; i < 2; i++) {
143         ret = compute_mask_matrix(mask_matrix[i], step_x[i], step_y[i]);
144         if (ret < 0)
145             return ret;
146     }
147     return ret;
148 }
149
150 int ff_opencl_apply_unsharp(AVFilterContext *ctx, AVFrame *in, AVFrame *out)
151 {
152     int ret;
153     AVFilterLink *link = ctx->inputs[0];
154     UnsharpContext *unsharp = ctx->priv;
155     cl_int status;
156     int cw = FF_CEIL_RSHIFT(link->w, unsharp->hsub);
157     int ch = FF_CEIL_RSHIFT(link->h, unsharp->vsub);
158     const size_t global_work_size = link->w * link->h + 2 * ch * cw;
159     FFOpenclParam opencl_param = {0};
160
161     opencl_param.ctx = ctx;
162     opencl_param.kernel = unsharp->opencl_ctx.kernel_env.kernel;
163     ret = ff_opencl_set_parameter(&opencl_param,
164                                   FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_inbuf),
165                                   FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_outbuf),
166                                   FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_luma_mask),
167                                   FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_chroma_mask),
168                                   FF_OPENCL_PARAM_INFO(unsharp->luma.amount),
169                                   FF_OPENCL_PARAM_INFO(unsharp->chroma.amount),
170                                   FF_OPENCL_PARAM_INFO(unsharp->luma.steps_x),
171                                   FF_OPENCL_PARAM_INFO(unsharp->luma.steps_y),
172                                   FF_OPENCL_PARAM_INFO(unsharp->chroma.steps_x),
173                                   FF_OPENCL_PARAM_INFO(unsharp->chroma.steps_y),
174                                   FF_OPENCL_PARAM_INFO(unsharp->luma.scalebits),
175                                   FF_OPENCL_PARAM_INFO(unsharp->chroma.scalebits),
176                                   FF_OPENCL_PARAM_INFO(unsharp->luma.halfscale),
177                                   FF_OPENCL_PARAM_INFO(unsharp->chroma.halfscale),
178                                   FF_OPENCL_PARAM_INFO(in->linesize[0]),
179                                   FF_OPENCL_PARAM_INFO(in->linesize[1]),
180                                   FF_OPENCL_PARAM_INFO(out->linesize[0]),
181                                   FF_OPENCL_PARAM_INFO(out->linesize[1]),
182                                   FF_OPENCL_PARAM_INFO(link->h),
183                                   FF_OPENCL_PARAM_INFO(link->w),
184                                   FF_OPENCL_PARAM_INFO(ch),
185                                   FF_OPENCL_PARAM_INFO(cw),
186                                   NULL);
187     if (ret < 0)
188         return ret;
189     status = clEnqueueNDRangeKernel(unsharp->opencl_ctx.kernel_env.command_queue,
190                                     unsharp->opencl_ctx.kernel_env.kernel, 1, NULL,
191                                     &global_work_size, NULL, 0, NULL, NULL);
192     if (status != CL_SUCCESS) {
193         av_log(ctx, AV_LOG_ERROR, "OpenCL run kernel error occurred: %s\n", av_opencl_errstr(status));
194         return AVERROR_EXTERNAL;
195     }
196     clFinish(unsharp->opencl_ctx.kernel_env.command_queue);
197     return av_opencl_buffer_read_image(out->data, unsharp->opencl_ctx.out_plane_size,
198                                        unsharp->opencl_ctx.plane_num, unsharp->opencl_ctx.cl_outbuf,
199                                        unsharp->opencl_ctx.cl_outbuf_size);
200 }
201
202 int ff_opencl_unsharp_init(AVFilterContext *ctx)
203 {
204     int ret = 0;
205     UnsharpContext *unsharp = ctx->priv;
206     ret = av_opencl_init(NULL);
207     if (ret < 0)
208         return ret;
209     ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_luma_mask,
210                                   sizeof(uint32_t) * (2 * unsharp->luma.steps_x + 1) * (2 * unsharp->luma.steps_y + 1),
211                                   CL_MEM_READ_ONLY, NULL);
212     if (ret < 0)
213         return ret;
214     ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_chroma_mask,
215                                   sizeof(uint32_t) * (2 * unsharp->chroma.steps_x + 1) * (2 * unsharp->chroma.steps_y + 1),
216                                   CL_MEM_READ_ONLY, NULL);
217     if (ret < 0)
218         return ret;
219     ret = generate_mask(ctx);
220     if (ret < 0)
221         return ret;
222     unsharp->opencl_ctx.plane_num = PLANE_NUM;
223     if (!unsharp->opencl_ctx.kernel_env.kernel) {
224         ret = av_opencl_create_kernel(&unsharp->opencl_ctx.kernel_env, "unsharp");
225         if (ret < 0) {
226             av_log(ctx, AV_LOG_ERROR, "OpenCL failed to create kernel with name 'unsharp'\n");
227             return ret;
228         }
229     }
230     return ret;
231 }
232
233 void ff_opencl_unsharp_uninit(AVFilterContext *ctx)
234 {
235     UnsharpContext *unsharp = ctx->priv;
236     av_opencl_buffer_release(&unsharp->opencl_ctx.cl_inbuf);
237     av_opencl_buffer_release(&unsharp->opencl_ctx.cl_outbuf);
238     av_opencl_buffer_release(&unsharp->opencl_ctx.cl_luma_mask);
239     av_opencl_buffer_release(&unsharp->opencl_ctx.cl_chroma_mask);
240     av_opencl_release_kernel(&unsharp->opencl_ctx.kernel_env);
241     av_opencl_uninit();
242 }
243
244 int ff_opencl_unsharp_process_inout_buf(AVFilterContext *ctx, AVFrame *in, AVFrame *out)
245 {
246     int ret = 0;
247     AVFilterLink *link = ctx->inputs[0];
248     UnsharpContext *unsharp = ctx->priv;
249     int ch = FF_CEIL_RSHIFT(link->h, unsharp->vsub);
250
251     if ((!unsharp->opencl_ctx.cl_inbuf) || (!unsharp->opencl_ctx.cl_outbuf)) {
252         unsharp->opencl_ctx.in_plane_size[0]  = (in->linesize[0] * in->height);
253         unsharp->opencl_ctx.in_plane_size[1]  = (in->linesize[1] * ch);
254         unsharp->opencl_ctx.in_plane_size[2]  = (in->linesize[2] * ch);
255         unsharp->opencl_ctx.out_plane_size[0] = (out->linesize[0] * out->height);
256         unsharp->opencl_ctx.out_plane_size[1] = (out->linesize[1] * ch);
257         unsharp->opencl_ctx.out_plane_size[2] = (out->linesize[2] * ch);
258         unsharp->opencl_ctx.cl_inbuf_size  = unsharp->opencl_ctx.in_plane_size[0] +
259                                              unsharp->opencl_ctx.in_plane_size[1] +
260                                              unsharp->opencl_ctx.in_plane_size[2];
261         unsharp->opencl_ctx.cl_outbuf_size = unsharp->opencl_ctx.out_plane_size[0] +
262                                              unsharp->opencl_ctx.out_plane_size[1] +
263                                              unsharp->opencl_ctx.out_plane_size[2];
264         if (!unsharp->opencl_ctx.cl_inbuf) {
265             ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_inbuf,
266                                           unsharp->opencl_ctx.cl_inbuf_size,
267                                           CL_MEM_READ_ONLY, NULL);
268             if (ret < 0)
269                 return ret;
270         }
271         if (!unsharp->opencl_ctx.cl_outbuf) {
272             ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_outbuf,
273                                           unsharp->opencl_ctx.cl_outbuf_size,
274                                           CL_MEM_READ_WRITE, NULL);
275             if (ret < 0)
276                 return ret;
277         }
278     }
279     return av_opencl_buffer_write_image(unsharp->opencl_ctx.cl_inbuf,
280                                         unsharp->opencl_ctx.cl_inbuf_size,
281                                         0, in->data, unsharp->opencl_ctx.in_plane_size,
282                                         unsharp->opencl_ctx.plane_num);
283 }