git.sesse.net Git - ffmpeg/blob - libavfilter/vf_unsharp_opencl.c

   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
  19 #include "libavutil/common.h"
  20 #include "libavutil/imgutils.h"
  21 #include "libavutil/mem.h"
  22 #include "libavutil/opt.h"
  23 #include "libavutil/pixdesc.h"
  24
  25 #include "avfilter.h"
  26 #include "internal.h"
  27 #include "opencl.h"
  28 #include "opencl_source.h"
  29 #include "video.h"
  30
  31 #define MAX_DIAMETER 23
  32
  33 typedef struct UnsharpOpenCLContext {
  34     OpenCLFilterContext ocf;
  35
  36     int              initialised;
  37     cl_kernel        kernel;
  38     cl_command_queue command_queue;
  39
  40     float luma_size_x;
  41     float luma_size_y;
  42     float luma_amount;
  43     float chroma_size_x;
  44     float chroma_size_y;
  45     float chroma_amount;
  46
  47     int global;
  48
  49     int nb_planes;
  50     struct {
  51         float blur_x[MAX_DIAMETER];
  52         float blur_y[MAX_DIAMETER];
  53
  54         cl_mem   matrix;
  55         cl_mem   coef_x;
  56         cl_mem   coef_y;
  57
  58         cl_int   size_x;
  59         cl_int   size_y;
  60         cl_float amount;
  61         cl_float threshold;
  62     } plane[4];
  63 } UnsharpOpenCLContext;
  64
  65
  66 static int unsharp_opencl_init(AVFilterContext *avctx)
  67 {
  68     UnsharpOpenCLContext *ctx = avctx->priv;
  69     cl_int cle;
  70     int err;
  71
  72     err = ff_opencl_filter_load_program(avctx, &ff_opencl_source_unsharp, 1);
  73     if (err < 0)
  74         goto fail;
  75
  76     ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
  77                                               ctx->ocf.hwctx->device_id,
  78                                               0, &cle);
  79     if (!ctx->command_queue) {
  80         av_log(avctx, AV_LOG_ERROR, "Failed to create OpenCL "
  81                "command queue: %d.\n", cle);
  82         err = AVERROR(EIO);
  83         goto fail;
  84     }
  85
  86     // Use global kernel if mask size will be too big for the local store..
  87     ctx->global = (ctx->luma_size_x   > 17.0f ||
  88                    ctx->luma_size_y   > 17.0f ||
  89                    ctx->chroma_size_x > 17.0f ||
  90                    ctx->chroma_size_y > 17.0f);
  91
  92     ctx->kernel = clCreateKernel(ctx->ocf.program,
  93                                  ctx->global ? "unsharp_global"
  94                                              : "unsharp_local", &cle);
  95     if (!ctx->kernel) {
  96         av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle);
  97         err = AVERROR(EIO);
  98         goto fail;
  99     }
 100
 101     ctx->initialised = 1;
 102     return 0;
 103
 104 fail:
 105     if (ctx->command_queue)
 106         clReleaseCommandQueue(ctx->command_queue);
 107     if (ctx->kernel)
 108         clReleaseKernel(ctx->kernel);
 109     return err;
 110 }
 111
 112 static int unsharp_opencl_make_filter_params(AVFilterContext *avctx)
 113 {
 114     UnsharpOpenCLContext *ctx = avctx->priv;
 115     const AVPixFmtDescriptor *desc;
 116     float *matrix;
 117     double val, sum;
 118     cl_int cle;
 119     cl_mem buffer;
 120     size_t matrix_bytes;
 121     float diam_x, diam_y, amount;
 122     int err, p, x, y, size_x, size_y;
 123
 124     desc = av_pix_fmt_desc_get(ctx->ocf.output_format);
 125
 126     ctx->nb_planes = 0;
 127     for (p = 0; p < desc->nb_components; p++)
 128         ctx->nb_planes = FFMAX(ctx->nb_planes, desc->comp[p].plane + 1);
 129
 130     for (p = 0; p < ctx->nb_planes; p++) {
 131         if (p == 0 || (desc->flags & AV_PIX_FMT_FLAG_RGB)) {
 132             diam_x = ctx->luma_size_x;
 133             diam_y = ctx->luma_size_y;
 134             amount = ctx->luma_amount;
 135         } else {
 136             diam_x = ctx->chroma_size_x;
 137             diam_y = ctx->chroma_size_y;
 138             amount = ctx->chroma_amount;
 139         }
 140         size_x = (int)ceil(diam_x) | 1;
 141         size_y = (int)ceil(diam_y) | 1;
 142         matrix_bytes = size_x * size_y * sizeof(float);
 143
 144         matrix = av_malloc(matrix_bytes);
 145         if (!matrix) {
 146             err = AVERROR(ENOMEM);
 147             goto fail;
 148         }
 149
 150         sum = 0.0;
 151         for (x = 0; x < size_x; x++) {
 152             double dx = (double)(x - size_x / 2) / diam_x;
 153             sum += ctx->plane[p].blur_x[x] = exp(-16.0 * (dx * dx));
 154         }
 155         for (x = 0; x < size_x; x++)
 156             ctx->plane[p].blur_x[x] /= sum;
 157
 158         sum = 0.0;
 159         for (y = 0; y < size_y; y++) {
 160             double dy = (double)(y - size_y / 2) / diam_y;
 161             sum += ctx->plane[p].blur_y[y] = exp(-16.0 * (dy * dy));
 162         }
 163         for (y = 0; y < size_y; y++)
 164             ctx->plane[p].blur_y[y] /= sum;
 165
 166         for (y = 0; y < size_y; y++) {
 167             for (x = 0; x < size_x; x++) {
 168                 val = ctx->plane[p].blur_x[x] * ctx->plane[p].blur_y[y];
 169                 matrix[y * size_x + x] = val;
 170             }
 171         }
 172
 173         if (ctx->global) {
 174             buffer = clCreateBuffer(ctx->ocf.hwctx->context,
 175                                     CL_MEM_READ_ONLY     |
 176                                     CL_MEM_COPY_HOST_PTR |
 177                                     CL_MEM_HOST_NO_ACCESS,
 178                                     matrix_bytes, matrix, &cle);
 179             if (!buffer) {
 180                 av_log(avctx, AV_LOG_ERROR, "Failed to create matrix buffer: "
 181                        "%d.\n", cle);
 182                 err = AVERROR(EIO);
 183                 goto fail;
 184             }
 185             ctx->plane[p].matrix = buffer;
 186         } else {
 187             buffer = clCreateBuffer(ctx->ocf.hwctx->context,
 188                                     CL_MEM_READ_ONLY     |
 189                                     CL_MEM_COPY_HOST_PTR |
 190                                     CL_MEM_HOST_NO_ACCESS,
 191                                     sizeof(ctx->plane[p].blur_x),
 192                                     ctx->plane[p].blur_x, &cle);
 193             if (!buffer) {
 194                 av_log(avctx, AV_LOG_ERROR, "Failed to create x-coef buffer: "
 195                        "%d.\n", cle);
 196                 err = AVERROR(EIO);
 197                 goto fail;
 198             }
 199             ctx->plane[p].coef_x = buffer;
 200
 201             buffer = clCreateBuffer(ctx->ocf.hwctx->context,
 202                                     CL_MEM_READ_ONLY     |
 203                                     CL_MEM_COPY_HOST_PTR |
 204                                     CL_MEM_HOST_NO_ACCESS,
 205                                     sizeof(ctx->plane[p].blur_y),
 206                                     ctx->plane[p].blur_y, &cle);
 207             if (!buffer) {
 208                 av_log(avctx, AV_LOG_ERROR, "Failed to create y-coef buffer: "
 209                        "%d.\n", cle);
 210                 err = AVERROR(EIO);
 211                 goto fail;
 212             }
 213             ctx->plane[p].coef_y = buffer;
 214         }
 215
 216         av_freep(&matrix);
 217
 218         ctx->plane[p].size_x = size_x;
 219         ctx->plane[p].size_y = size_y;
 220         ctx->plane[p].amount = amount;
 221     }
 222
 223     err = 0;
 224 fail:
 225     av_freep(&matrix);
 226     return err;
 227 }
 228
 229 static int unsharp_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
 230 {
 231     AVFilterContext    *avctx = inlink->dst;
 232     AVFilterLink     *outlink = avctx->outputs[0];
 233     UnsharpOpenCLContext *ctx = avctx->priv;
 234     AVFrame *output = NULL;
 235     cl_int cle;
 236     size_t global_work[2];
 237     size_t local_work[2];
 238     cl_mem src, dst;
 239     int err, p;
 240
 241     av_log(ctx, AV_LOG_DEBUG, "Filter input: %s, %ux%u (%"PRId64").\n",
 242            av_get_pix_fmt_name(input->format),
 243            input->width, input->height, input->pts);
 244
 245     if (!input->hw_frames_ctx)
 246         return AVERROR(EINVAL);
 247
 248     if (!ctx->initialised) {
 249         err = unsharp_opencl_init(avctx);
 250         if (err < 0)
 251             goto fail;
 252
 253         err = unsharp_opencl_make_filter_params(avctx);
 254         if (err < 0)
 255             goto fail;
 256     }
 257
 258     output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
 259     if (!output) {
 260         err = AVERROR(ENOMEM);
 261         goto fail;
 262     }
 263
 264     for (p = 0; p < FF_ARRAY_ELEMS(output->data); p++) {
 265         src = (cl_mem) input->data[p];
 266         dst = (cl_mem)output->data[p];
 267
 268         if (!dst)
 269             break;
 270
 271         CL_SET_KERNEL_ARG(ctx->kernel, 0, cl_mem,   &dst);
 272         CL_SET_KERNEL_ARG(ctx->kernel, 1, cl_mem,   &src);
 273         CL_SET_KERNEL_ARG(ctx->kernel, 2, cl_int,   &ctx->plane[p].size_x);
 274         CL_SET_KERNEL_ARG(ctx->kernel, 3, cl_int,   &ctx->plane[p].size_y);
 275         CL_SET_KERNEL_ARG(ctx->kernel, 4, cl_float, &ctx->plane[p].amount);
 276
 277         if (ctx->global) {
 278             CL_SET_KERNEL_ARG(ctx->kernel, 5, cl_mem, &ctx->plane[p].matrix);
 279         } else {
 280             CL_SET_KERNEL_ARG(ctx->kernel, 5, cl_mem, &ctx->plane[p].coef_x);
 281             CL_SET_KERNEL_ARG(ctx->kernel, 6, cl_mem, &ctx->plane[p].coef_y);
 282         }
 283
 284         err = ff_opencl_filter_work_size_from_image(avctx, global_work, output, p,
 285                                                     ctx->global ? 0 : 16);
 286         if (err < 0)
 287             goto fail;
 288
 289         local_work[0]  = 16;
 290         local_work[1]  = 16;
 291
 292         av_log(avctx, AV_LOG_DEBUG, "Run kernel on plane %d "
 293                "(%"SIZE_SPECIFIER"x%"SIZE_SPECIFIER").\n",
 294                p, global_work[0], global_work[1]);
 295
 296         cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
 297                                      global_work, ctx->global ? NULL : local_work,
 298                                      0, NULL, NULL);
 299         if (cle != CL_SUCCESS) {
 300             av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n",
 301                    cle);
 302             err = AVERROR(EIO);
 303             goto fail;
 304         }
 305     }
 306
 307     cle = clFinish(ctx->command_queue);
 308     if (cle != CL_SUCCESS) {
 309         av_log(avctx, AV_LOG_ERROR, "Failed to finish command queue: %d.\n",
 310                cle);
 311         err = AVERROR(EIO);
 312         goto fail;
 313     }
 314
 315     err = av_frame_copy_props(output, input);
 316     if (err < 0)
 317         goto fail;
 318
 319     av_frame_free(&input);
 320
 321     av_log(ctx, AV_LOG_DEBUG, "Filter output: %s, %ux%u (%"PRId64").\n",
 322            av_get_pix_fmt_name(output->format),
 323            output->width, output->height, output->pts);
 324
 325     return ff_filter_frame(outlink, output);
 326
 327 fail:
 328     clFinish(ctx->command_queue);
 329     av_frame_free(&input);
 330     av_frame_free(&output);
 331     return err;
 332 }
 333
 334 static av_cold void unsharp_opencl_uninit(AVFilterContext *avctx)
 335 {
 336     UnsharpOpenCLContext *ctx = avctx->priv;
 337     cl_int cle;
 338     int i;
 339
 340     for (i = 0; i < ctx->nb_planes; i++) {
 341         if (ctx->plane[i].matrix)
 342             clReleaseMemObject(ctx->plane[i].matrix);
 343         if (ctx->plane[i].coef_x)
 344             clReleaseMemObject(ctx->plane[i].coef_x);
 345         if (ctx->plane[i].coef_y)
 346             clReleaseMemObject(ctx->plane[i].coef_y);
 347     }
 348
 349     if (ctx->kernel) {
 350         cle = clReleaseKernel(ctx->kernel);
 351         if (cle != CL_SUCCESS)
 352             av_log(avctx, AV_LOG_ERROR, "Failed to release "
 353                    "kernel: %d.\n", cle);
 354     }
 355
 356     if (ctx->command_queue) {
 357         cle = clReleaseCommandQueue(ctx->command_queue);
 358         if (cle != CL_SUCCESS)
 359             av_log(avctx, AV_LOG_ERROR, "Failed to release "
 360                    "command queue: %d.\n", cle);
 361     }
 362
 363     ff_opencl_filter_uninit(avctx);
 364 }
 365
 366 #define OFFSET(x) offsetof(UnsharpOpenCLContext, x)
 367 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
 368 static const AVOption unsharp_opencl_options[] = {
 369     { "luma_msize_x",     "Set luma mask horizontal diameter (pixels)",
 370       OFFSET(luma_size_x),     AV_OPT_TYPE_FLOAT,
 371       { .dbl = 5.0 },   1, MAX_DIAMETER, FLAGS },
 372     { "lx",               "Set luma mask horizontal diameter (pixels)",
 373       OFFSET(luma_size_x),     AV_OPT_TYPE_FLOAT,
 374       { .dbl = 5.0 },   1, MAX_DIAMETER, FLAGS },
 375     { "luma_msize_y",     "Set luma mask vertical diameter (pixels)",
 376       OFFSET(luma_size_y),     AV_OPT_TYPE_FLOAT,
 377       { .dbl = 5.0 },   1, MAX_DIAMETER, FLAGS },
 378     { "ly",               "Set luma mask vertical diameter (pixels)",
 379       OFFSET(luma_size_y),     AV_OPT_TYPE_FLOAT,
 380       { .dbl = 5.0 },   1, MAX_DIAMETER, FLAGS },
 381     { "luma_amount",      "Set luma amount (multiplier)",
 382       OFFSET(luma_amount),     AV_OPT_TYPE_FLOAT,
 383       { .dbl = 1.0 }, -10, 10, FLAGS },
 384     { "la",               "Set luma amount (multiplier)",
 385       OFFSET(luma_amount),     AV_OPT_TYPE_FLOAT,
 386       { .dbl = 1.0 }, -10, 10, FLAGS },
 387
 388     { "chroma_msize_x",   "Set chroma mask horizontal diameter (pixels after subsampling)",
 389       OFFSET(chroma_size_x),   AV_OPT_TYPE_FLOAT,
 390       { .dbl = 5.0 },   1, MAX_DIAMETER, FLAGS },
 391     { "cx",               "Set chroma mask horizontal diameter (pixels after subsampling)",
 392       OFFSET(chroma_size_x),   AV_OPT_TYPE_FLOAT,
 393       { .dbl = 5.0 },   1, MAX_DIAMETER, FLAGS },
 394     { "chroma_msize_y",   "Set chroma mask vertical diameter (pixels after subsampling)",
 395       OFFSET(chroma_size_y),   AV_OPT_TYPE_FLOAT,
 396       { .dbl = 5.0 },   1, MAX_DIAMETER, FLAGS },
 397     { "cy",               "Set chroma mask vertical diameter (pixels after subsampling)",
 398       OFFSET(chroma_size_y),   AV_OPT_TYPE_FLOAT,
 399       { .dbl = 5.0 },   1, MAX_DIAMETER, FLAGS },
 400     { "chroma_amount",    "Set chroma amount (multiplier)",
 401       OFFSET(chroma_amount),   AV_OPT_TYPE_FLOAT,
 402       { .dbl = 0.0 }, -10, 10, FLAGS },
 403     { "ca",               "Set chroma amount (multiplier)",
 404       OFFSET(chroma_amount),   AV_OPT_TYPE_FLOAT,
 405       { .dbl = 0.0 }, -10, 10, FLAGS },
 406
 407     { NULL }
 408 };
 409
 410 AVFILTER_DEFINE_CLASS(unsharp_opencl);
 411
 412 static const AVFilterPad unsharp_opencl_inputs[] = {
 413     {
 414         .name         = "default",
 415         .type         = AVMEDIA_TYPE_VIDEO,
 416         .filter_frame = &unsharp_opencl_filter_frame,
 417         .config_props = &ff_opencl_filter_config_input,
 418     },
 419     { NULL }
 420 };
 421
 422 static const AVFilterPad unsharp_opencl_outputs[] = {
 423     {
 424         .name         = "default",
 425         .type         = AVMEDIA_TYPE_VIDEO,
 426         .config_props = &ff_opencl_filter_config_output,
 427     },
 428     { NULL }
 429 };
 430
 431 AVFilter ff_vf_unsharp_opencl = {
 432     .name           = "unsharp_opencl",
 433     .description    = NULL_IF_CONFIG_SMALL("Apply unsharp mask to input video"),
 434     .priv_size      = sizeof(UnsharpOpenCLContext),
 435     .priv_class     = &unsharp_opencl_class,
 436     .init           = &ff_opencl_filter_init,
 437     .uninit         = &unsharp_opencl_uninit,
 438     .query_formats  = &ff_opencl_filter_query_formats,
 439     .inputs         = unsharp_opencl_inputs,
 440     .outputs        = unsharp_opencl_outputs,
 441     .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
 442 };