Lines Matching defs:unsharp
24 * unsharp input video
131 UnsharpContext *unsharp = ctx->priv;
134 mask_matrix[0] = unsharp->opencl_ctx.cl_luma_mask;
135 mask_matrix[1] = unsharp->opencl_ctx.cl_chroma_mask;
136 step_x[0] = unsharp->luma.steps_x;
137 step_x[1] = unsharp->chroma.steps_x;
138 step_y[0] = unsharp->luma.steps_y;
139 step_y[1] = unsharp->chroma.steps_y;
143 unsharp->opencl_ctx.use_fast_kernels = 0;
145 unsharp->opencl_ctx.use_fast_kernels = 1;
163 UnsharpContext *unsharp = ctx->priv;
169 int cw = FF_CEIL_RSHIFT(link->w, unsharp->hsub);
170 int ch = FF_CEIL_RSHIFT(link->h, unsharp->vsub);
176 if (unsharp->opencl_ctx.use_fast_kernels) {
183 kernel1.kernel = unsharp->opencl_ctx.kernel_luma;
185 FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_inbuf),
186 FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_outbuf),
187 FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_luma_mask),
188 FF_OPENCL_PARAM_INFO(unsharp->luma.amount),
189 FF_OPENCL_PARAM_INFO(unsharp->luma.scalebits),
190 FF_OPENCL_PARAM_INFO(unsharp->luma.halfscale),
200 kernel2.kernel = unsharp->opencl_ctx.kernel_chroma;
202 FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_inbuf),
203 FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_outbuf),
204 FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_chroma_mask),
205 FF_OPENCL_PARAM_INFO(unsharp->chroma.amount),
206 FF_OPENCL_PARAM_INFO(unsharp->chroma.scalebits),
207 FF_OPENCL_PARAM_INFO(unsharp->chroma.halfscale),
219 status = clEnqueueNDRangeKernel(unsharp->opencl_ctx.command_queue,
220 unsharp->opencl_ctx.kernel_luma, 2, NULL,
222 status |=clEnqueueNDRangeKernel(unsharp->opencl_ctx.command_queue,
223 unsharp->opencl_ctx.kernel_chroma, 2, NULL,
231 kernel1.kernel = unsharp->opencl_ctx.kernel_default;
234 FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_inbuf),
235 FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_outbuf),
236 FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_luma_mask),
237 FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_chroma_mask),
238 FF_OPENCL_PARAM_INFO(unsharp->luma.amount),
239 FF_OPENCL_PARAM_INFO(unsharp->chroma.amount),
240 FF_OPENCL_PARAM_INFO(unsharp->luma.steps_x),
241 FF_OPENCL_PARAM_INFO(unsharp->luma.steps_y),
242 FF_OPENCL_PARAM_INFO(unsharp->chroma.steps_x),
243 FF_OPENCL_PARAM_INFO(unsharp->chroma.steps_y),
244 FF_OPENCL_PARAM_INFO(unsharp->luma.scalebits),
245 FF_OPENCL_PARAM_INFO(unsharp->chroma.scalebits),
246 FF_OPENCL_PARAM_INFO(unsharp->luma.halfscale),
247 FF_OPENCL_PARAM_INFO(unsharp->chroma.halfscale),
259 status = clEnqueueNDRangeKernel(unsharp->opencl_ctx.command_queue,
260 unsharp->opencl_ctx.kernel_default, 1, NULL,
267 clFinish(unsharp->opencl_ctx.command_queue);
268 return av_opencl_buffer_read_image(out->data, unsharp->opencl_ctx.out_plane_size,
269 unsharp->opencl_ctx.plane_num, unsharp->opencl_ctx.cl_outbuf,
270 unsharp->opencl_ctx.cl_outbuf_size);
277 UnsharpContext *unsharp = ctx->priv;
281 ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_luma_mask,
282 sizeof(uint32_t) * (2 * unsharp->luma.steps_x + 1) * (2 * unsharp->luma.steps_y + 1),
286 ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_chroma_mask,
287 sizeof(uint32_t) * (2 * unsharp->chroma.steps_x + 1) * (2 * unsharp->chroma.steps_y + 1),
294 unsharp->opencl_ctx.plane_num = PLANE_NUM;
295 unsharp->opencl_ctx.command_queue = av_opencl_get_command_queue();
296 if (!unsharp->opencl_ctx.command_queue) {
297 av_log(ctx, AV_LOG_ERROR, "Unable to get OpenCL command queue in filter 'unsharp'\n");
301 2*unsharp->luma.steps_x+1, 2*unsharp->luma.steps_y+1, 2*unsharp->chroma.steps_x+1, 2*unsharp->chroma.steps_y+1);
302 unsharp->opencl_ctx.program = av_opencl_compile("unsharp", build_opts);
303 if (!unsharp->opencl_ctx.program) {
304 av_log(ctx, AV_LOG_ERROR, "OpenCL failed to compile program 'unsharp'\n");
307 if (unsharp->opencl_ctx.use_fast_kernels) {
308 if (!unsharp->opencl_ctx.kernel_luma) {
309 unsharp->opencl_ctx.kernel_luma = clCreateKernel(unsharp->opencl_ctx.program, "unsharp_luma", &ret);
315 if (!unsharp->opencl_ctx.kernel_chroma) {
316 unsharp->opencl_ctx.kernel_chroma = clCreateKernel(unsharp->opencl_ctx.program, "unsharp_chroma", &ret);
324 if (!unsharp->opencl_ctx.kernel_default) {
325 unsharp->opencl_ctx.kernel_default = clCreateKernel(unsharp->opencl_ctx.program, "unsharp_default", &ret);
337 UnsharpContext *unsharp = ctx->priv;
338 av_opencl_buffer_release(&unsharp->opencl_ctx.cl_inbuf);
339 av_opencl_buffer_release(&unsharp->opencl_ctx.cl_outbuf);
340 av_opencl_buffer_release(&unsharp->opencl_ctx.cl_luma_mask);
341 av_opencl_buffer_release(&unsharp->opencl_ctx.cl_chroma_mask);
342 clReleaseKernel(unsharp->opencl_ctx.kernel_default);
343 clReleaseKernel(unsharp->opencl_ctx.kernel_luma);
344 clReleaseKernel(unsharp->opencl_ctx.kernel_chroma);
345 clReleaseProgram(unsharp->opencl_ctx.program);
346 unsharp->opencl_ctx.command_queue = NULL;
354 UnsharpContext *unsharp = ctx->priv;
355 int ch = FF_CEIL_RSHIFT(link->h, unsharp->vsub);
357 if ((!unsharp->opencl_ctx.cl_inbuf) || (!unsharp->opencl_ctx.cl_outbuf)) {
358 unsharp->opencl_ctx.in_plane_size[0] = (in->linesize[0] * in->height);
359 unsharp->opencl_ctx.in_plane_size[1] = (in->linesize[1] * ch);
360 unsharp->opencl_ctx.in_plane_size[2] = (in->linesize[2] * ch);
361 unsharp->opencl_ctx.out_plane_size[0] = (out->linesize[0] * out->height);
362 unsharp->opencl_ctx.out_plane_size[1] = (out->linesize[1] * ch);
363 unsharp->opencl_ctx.out_plane_size[2] = (out->linesize[2] * ch);
364 unsharp->opencl_ctx.cl_inbuf_size = unsharp->opencl_ctx.in_plane_size[0] +
365 unsharp->opencl_ctx.in_plane_size[1] +
366 unsharp->opencl_ctx.in_plane_size[2];
367 unsharp->opencl_ctx.cl_outbuf_size = unsharp->opencl_ctx.out_plane_size[0] +
368 unsharp->opencl_ctx.out_plane_size[1] +
369 unsharp->opencl_ctx.out_plane_size[2];
370 if (!unsharp->opencl_ctx.cl_inbuf) {
371 ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_inbuf,
372 unsharp->opencl_ctx.cl_inbuf_size,
377 if (!unsharp->opencl_ctx.cl_outbuf) {
378 ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_outbuf,
379 unsharp->opencl_ctx.cl_outbuf_size,
385 return av_opencl_buffer_write_image(unsharp->opencl_ctx.cl_inbuf,
386 unsharp->opencl_ctx.cl_inbuf_size,
387 0, in->data, unsharp->opencl_ctx.in_plane_size,
388 unsharp->opencl_ctx.plane_num);