Subversion Repositories Kolibri OS

Rev

Blame | Last modification | View Log | RSS feed

  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.  * transform input video
  24.  */
  25.  
  26. #include "libavutil/common.h"
  27. #include "libavutil/dict.h"
  28. #include "libavutil/pixdesc.h"
  29. #include "deshake_opencl.h"
  30. #include "libavutil/opencl_internal.h"
  31.  
  32. #define MATRIX_SIZE 6
  33. #define PLANE_NUM 3
  34.  
  35. int ff_opencl_transform(AVFilterContext *ctx,
  36.                         int width, int height, int cw, int ch,
  37.                         const float *matrix_y, const float *matrix_uv,
  38.                         enum InterpolateMethod interpolate,
  39.                         enum FillMethod fill, AVFrame *in, AVFrame *out)
  40. {
  41.     int ret = 0;
  42.     const size_t global_work_size = width * height + 2 * ch * cw;
  43.     cl_int status;
  44.     DeshakeContext *deshake = ctx->priv;
  45.     FFOpenclParam opencl_param = {0};
  46.  
  47.     opencl_param.ctx = ctx;
  48.     opencl_param.kernel = deshake->opencl_ctx.kernel_env.kernel;
  49.     ret = av_opencl_buffer_write(deshake->opencl_ctx.cl_matrix_y, (uint8_t *)matrix_y, deshake->opencl_ctx.matrix_size * sizeof(cl_float));
  50.     if (ret < 0)
  51.         return ret;
  52.     ret = av_opencl_buffer_write(deshake->opencl_ctx.cl_matrix_uv, (uint8_t *)matrix_uv, deshake->opencl_ctx.matrix_size * sizeof(cl_float));
  53.     if (ret < 0)
  54.         return ret;
  55.  
  56.     if ((unsigned int)interpolate > INTERPOLATE_BIQUADRATIC) {
  57.         av_log(ctx, AV_LOG_ERROR, "Selected interpolate method is invalid\n");
  58.         return AVERROR(EINVAL);
  59.     }
  60.     ret = ff_opencl_set_parameter(&opencl_param,
  61.                                   FF_OPENCL_PARAM_INFO(deshake->opencl_ctx.cl_inbuf),
  62.                                   FF_OPENCL_PARAM_INFO(deshake->opencl_ctx.cl_outbuf),
  63.                                   FF_OPENCL_PARAM_INFO(deshake->opencl_ctx.cl_matrix_y),
  64.                                   FF_OPENCL_PARAM_INFO(deshake->opencl_ctx.cl_matrix_uv),
  65.                                   FF_OPENCL_PARAM_INFO(interpolate),
  66.                                   FF_OPENCL_PARAM_INFO(fill),
  67.                                   FF_OPENCL_PARAM_INFO(in->linesize[0]),
  68.                                   FF_OPENCL_PARAM_INFO(out->linesize[0]),
  69.                                   FF_OPENCL_PARAM_INFO(in->linesize[1]),
  70.                                   FF_OPENCL_PARAM_INFO(out->linesize[1]),
  71.                                   FF_OPENCL_PARAM_INFO(height),
  72.                                   FF_OPENCL_PARAM_INFO(width),
  73.                                   FF_OPENCL_PARAM_INFO(ch),
  74.                                   FF_OPENCL_PARAM_INFO(cw),
  75.                                   NULL);
  76.     if (ret < 0)
  77.         return ret;
  78.     status = clEnqueueNDRangeKernel(deshake->opencl_ctx.kernel_env.command_queue,
  79.                                     deshake->opencl_ctx.kernel_env.kernel, 1, NULL,
  80.                                     &global_work_size, NULL, 0, NULL, NULL);
  81.     if (status != CL_SUCCESS) {
  82.         av_log(ctx, AV_LOG_ERROR, "OpenCL run kernel error occurred: %s\n", av_opencl_errstr(status));
  83.         return AVERROR_EXTERNAL;
  84.     }
  85.     clFinish(deshake->opencl_ctx.kernel_env.command_queue);
  86.     ret = av_opencl_buffer_read_image(out->data, deshake->opencl_ctx.out_plane_size,
  87.                                       deshake->opencl_ctx.plane_num, deshake->opencl_ctx.cl_outbuf,
  88.                                       deshake->opencl_ctx.cl_outbuf_size);
  89.     if (ret < 0)
  90.         return ret;
  91.     return ret;
  92. }
  93.  
  94. int ff_opencl_deshake_init(AVFilterContext *ctx)
  95. {
  96.     int ret = 0;
  97.     DeshakeContext *deshake = ctx->priv;
  98.     ret = av_opencl_init(NULL);
  99.     if (ret < 0)
  100.         return ret;
  101.     deshake->opencl_ctx.matrix_size = MATRIX_SIZE;
  102.     deshake->opencl_ctx.plane_num   = PLANE_NUM;
  103.     ret = av_opencl_buffer_create(&deshake->opencl_ctx.cl_matrix_y,
  104.         deshake->opencl_ctx.matrix_size*sizeof(cl_float), CL_MEM_READ_ONLY, NULL);
  105.     if (ret < 0)
  106.         return ret;
  107.     ret = av_opencl_buffer_create(&deshake->opencl_ctx.cl_matrix_uv,
  108.         deshake->opencl_ctx.matrix_size*sizeof(cl_float), CL_MEM_READ_ONLY, NULL);
  109.     if (ret < 0)
  110.         return ret;
  111.     if (!deshake->opencl_ctx.kernel_env.kernel) {
  112.         ret =  av_opencl_create_kernel(&deshake->opencl_ctx.kernel_env, "avfilter_transform");
  113.         if (ret < 0) {
  114.             av_log(ctx, AV_LOG_ERROR, "OpenCL failed to create kernel for name 'avfilter_transform'\n");
  115.             return ret;
  116.         }
  117.     }
  118.     return ret;
  119. }
  120.  
  121. void ff_opencl_deshake_uninit(AVFilterContext *ctx)
  122. {
  123.     DeshakeContext *deshake = ctx->priv;
  124.     av_opencl_buffer_release(&deshake->opencl_ctx.cl_inbuf);
  125.     av_opencl_buffer_release(&deshake->opencl_ctx.cl_outbuf);
  126.     av_opencl_buffer_release(&deshake->opencl_ctx.cl_matrix_y);
  127.     av_opencl_buffer_release(&deshake->opencl_ctx.cl_matrix_uv);
  128.     av_opencl_release_kernel(&deshake->opencl_ctx.kernel_env);
  129.     av_opencl_uninit();
  130. }
  131.  
  132.  
  133. int ff_opencl_deshake_process_inout_buf(AVFilterContext *ctx, AVFrame *in, AVFrame *out)
  134. {
  135.     int ret = 0;
  136.     AVFilterLink *link = ctx->inputs[0];
  137.     DeshakeContext *deshake = ctx->priv;
  138.     const int hshift = av_pix_fmt_desc_get(link->format)->log2_chroma_h;
  139.     int chroma_height = FF_CEIL_RSHIFT(link->h, hshift);
  140.  
  141.     if ((!deshake->opencl_ctx.cl_inbuf) || (!deshake->opencl_ctx.cl_outbuf)) {
  142.         deshake->opencl_ctx.in_plane_size[0]  = (in->linesize[0] * in->height);
  143.         deshake->opencl_ctx.in_plane_size[1]  = (in->linesize[1] * chroma_height);
  144.         deshake->opencl_ctx.in_plane_size[2]  = (in->linesize[2] * chroma_height);
  145.         deshake->opencl_ctx.out_plane_size[0] = (out->linesize[0] * out->height);
  146.         deshake->opencl_ctx.out_plane_size[1] = (out->linesize[1] * chroma_height);
  147.         deshake->opencl_ctx.out_plane_size[2] = (out->linesize[2] * chroma_height);
  148.         deshake->opencl_ctx.cl_inbuf_size  = deshake->opencl_ctx.in_plane_size[0] +
  149.                                              deshake->opencl_ctx.in_plane_size[1] +
  150.                                              deshake->opencl_ctx.in_plane_size[2];
  151.         deshake->opencl_ctx.cl_outbuf_size = deshake->opencl_ctx.out_plane_size[0] +
  152.                                              deshake->opencl_ctx.out_plane_size[1] +
  153.                                              deshake->opencl_ctx.out_plane_size[2];
  154.         if (!deshake->opencl_ctx.cl_inbuf) {
  155.             ret = av_opencl_buffer_create(&deshake->opencl_ctx.cl_inbuf,
  156.                                             deshake->opencl_ctx.cl_inbuf_size,
  157.                                             CL_MEM_READ_ONLY, NULL);
  158.             if (ret < 0)
  159.                 return ret;
  160.         }
  161.         if (!deshake->opencl_ctx.cl_outbuf) {
  162.             ret = av_opencl_buffer_create(&deshake->opencl_ctx.cl_outbuf,
  163.                                             deshake->opencl_ctx.cl_outbuf_size,
  164.                                             CL_MEM_READ_WRITE, NULL);
  165.             if (ret < 0)
  166.                 return ret;
  167.         }
  168.     }
  169.     ret = av_opencl_buffer_write_image(deshake->opencl_ctx.cl_inbuf,
  170.                                  deshake->opencl_ctx.cl_inbuf_size,
  171.                                  0, in->data,deshake->opencl_ctx.in_plane_size,
  172.                                  deshake->opencl_ctx.plane_num);
  173.     if(ret < 0)
  174.         return ret;
  175.     return ret;
  176. }
  177.