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         cle = clSetKernelArg(ctx->kernel, 0, sizeof(cl_mem), &dst);
 272         if (cle != CL_SUCCESS) {
 273             av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
 274                    "destination image argument: %d.\n", cle);
 275             goto fail;
 276         }
 277         cle = clSetKernelArg(ctx->kernel, 1, sizeof(cl_mem), &src);
 278         if (cle != CL_SUCCESS) {
 279             av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
 280                    "source image argument: %d.\n", cle);
 281             goto fail;
 282         }
 283         cle = clSetKernelArg(ctx->kernel, 2, sizeof(cl_int), &ctx->plane[p].size_x);
 284         if (cle != CL_SUCCESS) {
 285             av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
 286                    "matrix size argument: %d.\n", cle);
 287             goto fail;
 288         }
 289         cle = clSetKernelArg(ctx->kernel, 3, sizeof(cl_int), &ctx->plane[p].size_y);
 290         if (cle != CL_SUCCESS) {
 291             av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
 292                    "matrix size argument: %d.\n", cle);
 293             goto fail;
 294         }
 295         cle = clSetKernelArg(ctx->kernel, 4, sizeof(cl_float), &ctx->plane[p].amount);
 296         if (cle != CL_SUCCESS) {
 297             av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
 298                    "amount argument: %d.\n", cle);
 299             goto fail;
 300         }
 301         if (ctx->global) {
 302             cle = clSetKernelArg(ctx->kernel, 5, sizeof(cl_mem), &ctx->plane[p].matrix);
 303             if (cle != CL_SUCCESS) {
 304                 av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
 305                        "matrix argument: %d.\n", cle);
 306                 goto fail;
 307             }
 308         } else {
 309             cle = clSetKernelArg(ctx->kernel, 5, sizeof(cl_mem), &ctx->plane[p].coef_x);
 310             if (cle != CL_SUCCESS) {
 311                 av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
 312                        "x-coef argument: %d.\n", cle);
 313                 goto fail;
 314             }
 315             cle = clSetKernelArg(ctx->kernel, 6, sizeof(cl_mem), &ctx->plane[p].coef_y);
 316             if (cle != CL_SUCCESS) {
 317                 av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
 318                        "y-coef argument: %d.\n", cle);
 319                 goto fail;
 320             }
 321         }
 322
 323         if (ctx->global) {
 324             global_work[0] = output->width;
 325             global_work[1] = output->height;
 326         } else {
 327             global_work[0] = FFALIGN(output->width,  16);
 328             global_work[1] = FFALIGN(output->height, 16);
 329             local_work[0]  = 16;
 330             local_work[1]  = 16;
 331         }
 332
 333         av_log(avctx, AV_LOG_DEBUG, "Run kernel on plane %d "
 334                "(%zux%zu).\n", p, global_work[0], global_work[1]);
 335
 336         cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
 337                                      global_work, ctx->global ? NULL : local_work,
 338                                      0, NULL, NULL);
 339         if (cle != CL_SUCCESS) {
 340             av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n",
 341                    cle);
 342             err = AVERROR(EIO);
 343             goto fail;
 344         }
 345     }
 346
 347     cle = clFinish(ctx->command_queue);
 348     if (cle != CL_SUCCESS) {
 349         av_log(avctx, AV_LOG_ERROR, "Failed to finish command queue: %d.\n",
 350                cle);
 351         err = AVERROR(EIO);
 352         goto fail;
 353     }
 354
 355     err = av_frame_copy_props(output, input);
 356     if (err < 0)
 357         goto fail;
 358
 359     av_frame_free(&input);
 360
 361     av_log(ctx, AV_LOG_DEBUG, "Filter output: %s, %ux%u (%"PRId64").\n",
 362            av_get_pix_fmt_name(output->format),
 363            output->width, output->height, output->pts);
 364
 365     return ff_filter_frame(outlink, output);
 366
 367 fail:
 368     clFinish(ctx->command_queue);
 369     av_frame_free(&input);
 370     av_frame_free(&output);
 371     return err;
 372 }
 373
 374 static av_cold void unsharp_opencl_uninit(AVFilterContext *avctx)
 375 {
 376     UnsharpOpenCLContext *ctx = avctx->priv;
 377     cl_int cle;
 378     int i;
 379
 380     for (i = 0; i < ctx->nb_planes; i++) {
 381         if (ctx->plane[i].matrix)
 382             clReleaseMemObject(ctx->plane[i].matrix);
 383         if (ctx->plane[i].coef_x)
 384             clReleaseMemObject(ctx->plane[i].coef_x);
 385         if (ctx->plane[i].coef_y)
 386             clReleaseMemObject(ctx->plane[i].coef_y);
 387     }
 388
 389     if (ctx->kernel) {
 390         cle = clReleaseKernel(ctx->kernel);
 391         if (cle != CL_SUCCESS)
 392             av_log(avctx, AV_LOG_ERROR, "Failed to release "
 393                    "kernel: %d.\n", cle);
 394     }
 395
 396     if (ctx->command_queue) {
 397         cle = clReleaseCommandQueue(ctx->command_queue);
 398         if (cle != CL_SUCCESS)
 399             av_log(avctx, AV_LOG_ERROR, "Failed to release "
 400                    "command queue: %d.\n", cle);
 401     }
 402
 403     ff_opencl_filter_uninit(avctx);
 404 }
 405
 406 #define OFFSET(x) offsetof(UnsharpOpenCLContext, x)
 407 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
 408 static const AVOption unsharp_opencl_options[] = {
 409     { "luma_msize_x",     "Set luma mask horizontal diameter (pixels)",
 410       OFFSET(luma_size_x),     AV_OPT_TYPE_FLOAT,
 411       { .dbl = 5.0 },   1, MAX_DIAMETER, FLAGS },
 412     { "lx",               "Set luma mask horizontal diameter (pixels)",
 413       OFFSET(luma_size_x),     AV_OPT_TYPE_FLOAT,
 414       { .dbl = 5.0 },   1, MAX_DIAMETER, FLAGS },
 415     { "luma_msize_y",     "Set luma mask vertical diameter (pixels)",
 416       OFFSET(luma_size_y),     AV_OPT_TYPE_FLOAT,
 417       { .dbl = 5.0 },   1, MAX_DIAMETER, FLAGS },
 418     { "ly",               "Set luma mask vertical diameter (pixels)",
 419       OFFSET(luma_size_y),     AV_OPT_TYPE_FLOAT,
 420       { .dbl = 5.0 },   1, MAX_DIAMETER, FLAGS },
 421     { "luma_amount",      "Set luma amount (multiplier)",
 422       OFFSET(luma_amount),     AV_OPT_TYPE_FLOAT,
 423       { .dbl = 1.0 }, -10, 10, FLAGS },
 424     { "la",               "Set luma amount (multiplier)",
 425       OFFSET(luma_amount),     AV_OPT_TYPE_FLOAT,
 426       { .dbl = 1.0 }, -10, 10, FLAGS },
 427
 428     { "chroma_msize_x",   "Set chroma mask horizontal diameter (pixels after subsampling)",
 429       OFFSET(chroma_size_x),   AV_OPT_TYPE_FLOAT,
 430       { .dbl = 5.0 },   1, MAX_DIAMETER, FLAGS },
 431     { "cx",               "Set chroma mask horizontal diameter (pixels after subsampling)",
 432       OFFSET(chroma_size_x),   AV_OPT_TYPE_FLOAT,
 433       { .dbl = 5.0 },   1, MAX_DIAMETER, FLAGS },
 434     { "chroma_msize_y",   "Set chroma mask vertical diameter (pixels after subsampling)",
 435       OFFSET(chroma_size_y),   AV_OPT_TYPE_FLOAT,
 436       { .dbl = 5.0 },   1, MAX_DIAMETER, FLAGS },
 437     { "cy",               "Set chroma mask vertical diameter (pixels after subsampling)",
 438       OFFSET(chroma_size_y),   AV_OPT_TYPE_FLOAT,
 439       { .dbl = 5.0 },   1, MAX_DIAMETER, FLAGS },
 440     { "chroma_amount",    "Set chroma amount (multiplier)",
 441       OFFSET(chroma_amount),   AV_OPT_TYPE_FLOAT,
 442       { .dbl = 0.0 }, -10, 10, FLAGS },
 443     { "ca",               "Set chroma amount (multiplier)",
 444       OFFSET(chroma_amount),   AV_OPT_TYPE_FLOAT,
 445       { .dbl = 0.0 }, -10, 10, FLAGS },
 446
 447     { NULL }
 448 };
 449
 450 AVFILTER_DEFINE_CLASS(unsharp_opencl);
 451
 452 static const AVFilterPad unsharp_opencl_inputs[] = {
 453     {
 454         .name         = "default",
 455         .type         = AVMEDIA_TYPE_VIDEO,
 456         .filter_frame = &unsharp_opencl_filter_frame,
 457         .config_props = &ff_opencl_filter_config_input,
 458     },
 459     { NULL }
 460 };
 461
 462 static const AVFilterPad unsharp_opencl_outputs[] = {
 463     {
 464         .name         = "default",
 465         .type         = AVMEDIA_TYPE_VIDEO,
 466         .config_props = &ff_opencl_filter_config_output,
 467     },
 468     { NULL }
 469 };
 470
 471 AVFilter ff_vf_unsharp_opencl = {
 472     .name           = "unsharp_opencl",
 473     .description    = NULL_IF_CONFIG_SMALL("Apply unsharp mask to input video"),
 474     .priv_size      = sizeof(UnsharpOpenCLContext),
 475     .priv_class     = &unsharp_opencl_class,
 476     .init           = &ff_opencl_filter_init,
 477     .uninit         = &unsharp_opencl_uninit,
 478     .query_formats  = &ff_opencl_filter_query_formats,
 479     .inputs         = unsharp_opencl_inputs,
 480     .outputs        = unsharp_opencl_outputs,
 481     .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
 482 };