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