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

   1 /*
   2  * Copyright (c) 2018 Danil Iashchenko
   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 #include "libavutil/common.h"
  22 #include "libavutil/imgutils.h"
  23 #include "libavutil/mem.h"
  24 #include "libavutil/opt.h"
  25 #include "libavutil/pixdesc.h"
  26 #include "libavutil/avstring.h"
  27
  28
  29 #include "avfilter.h"
  30 #include "internal.h"
  31 #include "opencl.h"
  32 #include "opencl_source.h"
  33 #include "video.h"
  34
  35 typedef struct NeighborOpenCLContext {
  36     OpenCLFilterContext ocf;
  37
  38     int              initialised;
  39     cl_kernel        kernel;
  40     cl_command_queue command_queue;
  41
  42     char *matrix_str[4];
  43
  44     cl_float threshold[4];
  45     cl_int coordinates;
  46     cl_mem coord;
  47
  48 } NeighborOpenCLContext;
  49
  50 static int neighbor_opencl_init(AVFilterContext *avctx)
  51 {
  52     NeighborOpenCLContext *ctx = avctx->priv;
  53     const char *kernel_name;
  54     cl_int cle;
  55     int err;
  56
  57     err = ff_opencl_filter_load_program(avctx, &ff_opencl_source_neighbor, 1);
  58     if (err < 0)
  59         goto fail;
  60
  61     ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
  62                                               ctx->ocf.hwctx->device_id,
  63                                               0, &cle);
  64     CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL "
  65                      "command queue %d.\n", cle);
  66
  67     if (!strcmp(avctx->filter->name, "erosion_opencl")){
  68         kernel_name = "erosion_global";
  69     } else if (!strcmp(avctx->filter->name, "dilation_opencl")){
  70         kernel_name = "dilation_global";
  71     }
  72     ctx->kernel = clCreateKernel(ctx->ocf.program, kernel_name, &cle);
  73     CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create "
  74                      "kernel %d.\n", cle);
  75
  76     ctx->initialised = 1;
  77     return 0;
  78
  79 fail:
  80     if (ctx->command_queue)
  81         clReleaseCommandQueue(ctx->command_queue);
  82     if (ctx->kernel)
  83         clReleaseKernel(ctx->kernel);
  84     return err;
  85 }
  86
  87 static int neighbor_opencl_make_filter_params(AVFilterContext *avctx)
  88 {
  89     NeighborOpenCLContext *ctx = avctx->priv;
  90     cl_int matrix[9];
  91     cl_mem buffer;
  92     cl_int cle;
  93     int i;
  94
  95     for (i = 0; i < 4; i++) {
  96         ctx->threshold[i] /= 255.0;
  97     }
  98
  99     matrix[4] = 0;
 100     for (i = 0; i < 8; i++) {
 101         if (ctx->coordinates & (1 << i)) {
 102             matrix[i > 3 ? i + 1: i] = 1;
 103         }
 104     }
 105     buffer = clCreateBuffer(ctx->ocf.hwctx->context,
 106                             CL_MEM_READ_ONLY |
 107                             CL_MEM_COPY_HOST_PTR |
 108                             CL_MEM_HOST_NO_ACCESS,
 109                             9 * sizeof(cl_int), matrix, &cle);
 110     if (!buffer) {
 111         av_log(avctx, AV_LOG_ERROR, "Failed to create matrix buffer: "
 112                "%d.\n", cle);
 113         return AVERROR(EIO);
 114     }
 115     ctx->coord = buffer;
 116
 117     return 0;
 118 }
 119
 120
 121 static int neighbor_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
 122 {
 123     AVFilterContext *avctx = inlink->dst;
 124     AVFilterLink *outlink = avctx->outputs[0];
 125     NeighborOpenCLContext *ctx = avctx->priv;
 126     AVFrame *output = NULL;
 127     cl_int cle;
 128     size_t global_work[2];
 129     cl_mem src, dst;
 130     int err, p;
 131     size_t origin[3] = {0, 0, 0};
 132     size_t region[3] = {0, 0, 1};
 133
 134     av_log(ctx, AV_LOG_DEBUG, "Filter input: %s, %ux%u (%"PRId64").\n",
 135            av_get_pix_fmt_name(input->format),
 136            input->width, input->height, input->pts);
 137
 138     if (!input->hw_frames_ctx)
 139         return AVERROR(EINVAL);
 140
 141     if (!ctx->initialised) {
 142         err = neighbor_opencl_init(avctx);
 143         if (err < 0)
 144             goto fail;
 145
 146         err = neighbor_opencl_make_filter_params(avctx);
 147         if (err < 0)
 148             goto fail;
 149
 150     }
 151
 152     output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
 153     if (!output) {
 154         err = AVERROR(ENOMEM);
 155         goto fail;
 156     }
 157
 158     for (p = 0; p < FF_ARRAY_ELEMS(output->data); p++) {
 159         src = (cl_mem) input->data[p];
 160         dst = (cl_mem)output->data[p];
 161
 162         if (!dst)
 163             break;
 164
 165         if (ctx->threshold[p] == 0) {
 166             err = ff_opencl_filter_work_size_from_image(avctx, region, output, p, 0);
 167             if (err < 0)
 168                 goto fail;
 169
 170             cle = clEnqueueCopyImage(ctx->command_queue, src, dst,
 171                                      origin, origin, region, 0, NULL, NULL);
 172             CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to copy plane %d: %d.\n",
 173                              p, cle);
 174         } else {
 175             CL_SET_KERNEL_ARG(ctx->kernel, 0, cl_mem,   &dst);
 176             CL_SET_KERNEL_ARG(ctx->kernel, 1, cl_mem,   &src);
 177             CL_SET_KERNEL_ARG(ctx->kernel, 2, cl_float, &ctx->threshold[p]);
 178             CL_SET_KERNEL_ARG(ctx->kernel, 3, cl_mem,   &ctx->coord);
 179
 180             err = ff_opencl_filter_work_size_from_image(avctx, global_work, output, p, 0);
 181             if (err < 0)
 182                 goto fail;
 183
 184             av_log(avctx, AV_LOG_DEBUG, "Run kernel on plane %d "
 185                    "(%"SIZE_SPECIFIER"x%"SIZE_SPECIFIER").\n",
 186                    p, global_work[0], global_work[1]);
 187
 188             cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
 189                                          global_work, NULL,
 190                                          0, NULL, NULL);
 191             CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue "
 192                              "kernel: %d.\n", cle);
 193         }
 194     }
 195
 196     cle = clFinish(ctx->command_queue);
 197     CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
 198
 199     err = av_frame_copy_props(output, input);
 200     if (err < 0)
 201         goto fail;
 202
 203     av_frame_free(&input);
 204
 205     av_log(ctx, AV_LOG_DEBUG, "Filter output: %s, %ux%u (%"PRId64").\n",
 206            av_get_pix_fmt_name(output->format),
 207            output->width, output->height, output->pts);
 208
 209     return ff_filter_frame(outlink, output);
 210
 211 fail:
 212     clFinish(ctx->command_queue);
 213     av_frame_free(&input);
 214     av_frame_free(&output);
 215     return err;
 216 }
 217
 218 static av_cold void neighbor_opencl_uninit(AVFilterContext *avctx)
 219 {
 220     NeighborOpenCLContext *ctx = avctx->priv;
 221     cl_int cle;
 222
 223     clReleaseMemObject(ctx->coord);
 224
 225     if (ctx->kernel) {
 226         cle = clReleaseKernel(ctx->kernel);
 227         if (cle != CL_SUCCESS)
 228             av_log(avctx, AV_LOG_ERROR, "Failed to release "
 229                    "kernel: %d.\n", cle);
 230     }
 231
 232     if (ctx->command_queue) {
 233         cle = clReleaseCommandQueue(ctx->command_queue);
 234         if (cle != CL_SUCCESS)
 235             av_log(avctx, AV_LOG_ERROR, "Failed to release "
 236                    "command queue: %d.\n", cle);
 237     }
 238
 239     ff_opencl_filter_uninit(avctx);
 240 }
 241
 242 static const AVFilterPad neighbor_opencl_inputs[] = {
 243     {
 244         .name         = "default",
 245         .type         = AVMEDIA_TYPE_VIDEO,
 246         .filter_frame = &neighbor_opencl_filter_frame,
 247         .config_props = &ff_opencl_filter_config_input,
 248     },
 249     { NULL }
 250 };
 251
 252 static const AVFilterPad neighbor_opencl_outputs[] = {
 253     {
 254         .name         = "default",
 255         .type         = AVMEDIA_TYPE_VIDEO,
 256         .config_props = &ff_opencl_filter_config_output,
 257     },
 258     { NULL }
 259 };
 260
 261 #define OFFSET(x) offsetof(NeighborOpenCLContext, x)
 262 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
 263
 264 #if CONFIG_EROSION_OPENCL_FILTER
 265
 266 static const AVOption erosion_opencl_options[] = {
 267     { "threshold0",  "set threshold for 1st plane",   OFFSET(threshold[0]),   AV_OPT_TYPE_FLOAT, {.dbl=65535.0}, 0.0, 65535, FLAGS },
 268     { "threshold1",  "set threshold for 2nd plane",   OFFSET(threshold[1]),   AV_OPT_TYPE_FLOAT, {.dbl=65535.0}, 0.0, 65535, FLAGS },
 269     { "threshold2",  "set threshold for 3rd plane",   OFFSET(threshold[2]),   AV_OPT_TYPE_FLOAT, {.dbl=65535.0}, 0.0, 65535, FLAGS },
 270     { "threshold3",  "set threshold for 4th plane",   OFFSET(threshold[3]),   AV_OPT_TYPE_FLOAT, {.dbl=65535.0}, 0.0, 65535, FLAGS },
 271     { "coordinates", "set coordinates",               OFFSET(coordinates),    AV_OPT_TYPE_INT,   {.i64=255},     0,   255,   FLAGS },
 272     { NULL }
 273 };
 274
 275 AVFILTER_DEFINE_CLASS(erosion_opencl);
 276
 277 const AVFilter ff_vf_erosion_opencl = {
 278     .name           = "erosion_opencl",
 279     .description    = NULL_IF_CONFIG_SMALL("Apply erosion effect"),
 280     .priv_size      = sizeof(NeighborOpenCLContext),
 281     .priv_class     = &erosion_opencl_class,
 282     .init           = &ff_opencl_filter_init,
 283     .uninit         = &neighbor_opencl_uninit,
 284     .query_formats  = &ff_opencl_filter_query_formats,
 285     .inputs         = neighbor_opencl_inputs,
 286     .outputs        = neighbor_opencl_outputs,
 287     .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
 288 };
 289
 290 #endif /* CONFIG_EROSION_OPENCL_FILTER */
 291
 292 #if CONFIG_DILATION_OPENCL_FILTER
 293
 294 static const AVOption dilation_opencl_options[] = {
 295     { "threshold0",  "set threshold for 1st plane",   OFFSET(threshold[0]),   AV_OPT_TYPE_FLOAT, {.dbl=65535.0}, 0.0, 65535, FLAGS },
 296     { "threshold1",  "set threshold for 2nd plane",   OFFSET(threshold[1]),   AV_OPT_TYPE_FLOAT, {.dbl=65535.0}, 0.0, 65535, FLAGS },
 297     { "threshold2",  "set threshold for 3rd plane",   OFFSET(threshold[2]),   AV_OPT_TYPE_FLOAT, {.dbl=65535.0}, 0.0, 65535, FLAGS },
 298     { "threshold3",  "set threshold for 4th plane",   OFFSET(threshold[3]),   AV_OPT_TYPE_FLOAT, {.dbl=65535.0}, 0.0, 65535, FLAGS },
 299     { "coordinates", "set coordinates",               OFFSET(coordinates),    AV_OPT_TYPE_INT,   {.i64=255},     0,   255,   FLAGS },
 300     { NULL }
 301 };
 302
 303 AVFILTER_DEFINE_CLASS(dilation_opencl);
 304
 305 const AVFilter ff_vf_dilation_opencl = {
 306     .name           = "dilation_opencl",
 307     .description    = NULL_IF_CONFIG_SMALL("Apply dilation effect"),
 308     .priv_size      = sizeof(NeighborOpenCLContext),
 309     .priv_class     = &dilation_opencl_class,
 310     .init           = &ff_opencl_filter_init,
 311     .uninit         = &neighbor_opencl_uninit,
 312     .query_formats  = &ff_opencl_filter_query_formats,
 313     .inputs         = neighbor_opencl_inputs,
 314     .outputs        = neighbor_opencl_outputs,
 315     .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
 316 };
 317
 318 #endif /* CONFIG_DILATION_OPENCL_FILTER */