deshake_opencl.c
Go to the documentation of this file.
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 
31 #define MATRIX_SIZE 6
32 #define PLANE_NUM 3
33 
34 #define TRANSFORM_OPENCL_CHECK(method, ...) \
35  status = method(__VA_ARGS__); \
36  if (status != CL_SUCCESS) { \
37  av_log(ctx, AV_LOG_ERROR, "error %s %d\n", # method, status); \
38  return AVERROR_EXTERNAL; \
39  }
40 
41 #define TRANSFORM_OPENCL_SET_KERNEL_ARG(arg_ptr) \
42  status = clSetKernelArg((kernel),(arg_no++),(sizeof(arg_ptr)),(void*)(&(arg_ptr))); \
43  if (status != CL_SUCCESS) { \
44  av_log(ctx, AV_LOG_ERROR, "cannot set kernel argument: %d\n", status ); \
45  return AVERROR_EXTERNAL; \
46  }
47 
48 int ff_opencl_transform(AVFilterContext *ctx,
49  int width, int height, int cw, int ch,
50  const float *matrix_y, const float *matrix_uv,
51  enum InterpolateMethod interpolate,
52  enum FillMethod fill, AVFrame *in, AVFrame *out)
53 {
54  int arg_no, ret = 0;
55  const size_t global_work_size = width * height + 2 * ch * cw;
56  cl_kernel kernel;
57  cl_int status;
58  DeshakeContext *deshake = ctx->priv;
59  ret = av_opencl_buffer_write(deshake->opencl_ctx.cl_matrix_y, (uint8_t *)matrix_y, deshake->opencl_ctx.matrix_size * sizeof(cl_float));
60  if (ret < 0)
61  return ret;
62  ret = av_opencl_buffer_write(deshake->opencl_ctx.cl_matrix_uv, (uint8_t *)matrix_uv, deshake->opencl_ctx.matrix_size * sizeof(cl_float));
63  if (ret < 0)
64  return ret;
65  kernel = deshake->opencl_ctx.kernel_env.kernel;
66  arg_no = 0;
67 
68  if ((unsigned int)interpolate > INTERPOLATE_BIQUADRATIC) {
69  av_log(ctx, AV_LOG_ERROR, "Selected interpolate method is invalid\n");
70  return AVERROR(EINVAL);
71  }
72  TRANSFORM_OPENCL_SET_KERNEL_ARG(deshake->opencl_ctx.cl_inbuf);
73  TRANSFORM_OPENCL_SET_KERNEL_ARG(deshake->opencl_ctx.cl_outbuf);
74  TRANSFORM_OPENCL_SET_KERNEL_ARG(deshake->opencl_ctx.cl_matrix_y);
75  TRANSFORM_OPENCL_SET_KERNEL_ARG(deshake->opencl_ctx.cl_matrix_uv);
76  TRANSFORM_OPENCL_SET_KERNEL_ARG(interpolate);
77  TRANSFORM_OPENCL_SET_KERNEL_ARG(fill);
78  TRANSFORM_OPENCL_SET_KERNEL_ARG(in->linesize[0]);
79  TRANSFORM_OPENCL_SET_KERNEL_ARG(out->linesize[0]);
80  TRANSFORM_OPENCL_SET_KERNEL_ARG(in->linesize[1]);
81  TRANSFORM_OPENCL_SET_KERNEL_ARG(out->linesize[1]);
82  TRANSFORM_OPENCL_SET_KERNEL_ARG(height);
83  TRANSFORM_OPENCL_SET_KERNEL_ARG(width);
84  TRANSFORM_OPENCL_SET_KERNEL_ARG(ch);
85  TRANSFORM_OPENCL_SET_KERNEL_ARG(cw);
86  TRANSFORM_OPENCL_CHECK(clEnqueueNDRangeKernel, deshake->opencl_ctx.kernel_env.command_queue, deshake->opencl_ctx.kernel_env.kernel, 1, NULL,
87  &global_work_size, NULL, 0, NULL, NULL);
88  clFinish(deshake->opencl_ctx.kernel_env.command_queue);
89  ret = av_opencl_buffer_read_image(out->data, deshake->opencl_ctx.out_plane_size,
90  deshake->opencl_ctx.plane_num, deshake->opencl_ctx.cl_outbuf,
91  deshake->opencl_ctx.cl_outbuf_size);
92  if (ret < 0)
93  return ret;
94  return ret;
95 }
96 
97 int ff_opencl_deshake_init(AVFilterContext *ctx)
98 {
99  int ret = 0;
100  DeshakeContext *deshake = ctx->priv;
101  ret = av_opencl_init(NULL);
102  if (ret < 0)
103  return ret;
104  deshake->opencl_ctx.matrix_size = MATRIX_SIZE;
105  deshake->opencl_ctx.plane_num = PLANE_NUM;
106  ret = av_opencl_buffer_create(&deshake->opencl_ctx.cl_matrix_y,
107  deshake->opencl_ctx.matrix_size*sizeof(cl_float), CL_MEM_READ_ONLY, NULL);
108  if (ret < 0)
109  return ret;
110  ret = av_opencl_buffer_create(&deshake->opencl_ctx.cl_matrix_uv,
111  deshake->opencl_ctx.matrix_size*sizeof(cl_float), CL_MEM_READ_ONLY, NULL);
112  if (ret < 0)
113  return ret;
114  if (!deshake->opencl_ctx.kernel_env.kernel) {
115  ret = av_opencl_create_kernel(&deshake->opencl_ctx.kernel_env, "avfilter_transform");
116  if (ret < 0) {
117  av_log(ctx, AV_LOG_ERROR, "OpenCL failed to create kernel for name 'avfilter_transform'\n");
118  return ret;
119  }
120  }
121  return ret;
122 }
123 
124 void ff_opencl_deshake_uninit(AVFilterContext *ctx)
125 {
126  DeshakeContext *deshake = ctx->priv;
127  av_opencl_buffer_release(&deshake->opencl_ctx.cl_inbuf);
128  av_opencl_buffer_release(&deshake->opencl_ctx.cl_outbuf);
129  av_opencl_buffer_release(&deshake->opencl_ctx.cl_matrix_y);
130  av_opencl_buffer_release(&deshake->opencl_ctx.cl_matrix_uv);
131  av_opencl_release_kernel(&deshake->opencl_ctx.kernel_env);
132  av_opencl_uninit();
133 }
134 
135 
136 int ff_opencl_deshake_process_inout_buf(AVFilterContext *ctx, AVFrame *in, AVFrame *out)
137 {
138  int ret = 0;
139  AVFilterLink *link = ctx->inputs[0];
140  DeshakeContext *deshake = ctx->priv;
141  int chroma_height = -((-link->h) >> av_pix_fmt_desc_get(link->format)->log2_chroma_h);
142 
143  if ((!deshake->opencl_ctx.cl_inbuf) || (!deshake->opencl_ctx.cl_outbuf)) {
144  deshake->opencl_ctx.in_plane_size[0] = (in->linesize[0] * in->height);
145  deshake->opencl_ctx.in_plane_size[1] = (in->linesize[1] * chroma_height);
146  deshake->opencl_ctx.in_plane_size[2] = (in->linesize[2] * chroma_height);
147  deshake->opencl_ctx.out_plane_size[0] = (out->linesize[0] * out->height);
148  deshake->opencl_ctx.out_plane_size[1] = (out->linesize[1] * chroma_height);
149  deshake->opencl_ctx.out_plane_size[2] = (out->linesize[2] * chroma_height);
150  deshake->opencl_ctx.cl_inbuf_size = deshake->opencl_ctx.in_plane_size[0] +
151  deshake->opencl_ctx.in_plane_size[1] +
152  deshake->opencl_ctx.in_plane_size[2];
153  deshake->opencl_ctx.cl_outbuf_size = deshake->opencl_ctx.out_plane_size[0] +
154  deshake->opencl_ctx.out_plane_size[1] +
155  deshake->opencl_ctx.out_plane_size[2];
156  if (!deshake->opencl_ctx.cl_inbuf) {
157  ret = av_opencl_buffer_create(&deshake->opencl_ctx.cl_inbuf,
158  deshake->opencl_ctx.cl_inbuf_size,
159  CL_MEM_READ_ONLY, NULL);
160  if (ret < 0)
161  return ret;
162  }
163  if (!deshake->opencl_ctx.cl_outbuf) {
164  ret = av_opencl_buffer_create(&deshake->opencl_ctx.cl_outbuf,
165  deshake->opencl_ctx.cl_outbuf_size,
166  CL_MEM_READ_WRITE, NULL);
167  if (ret < 0)
168  return ret;
169  }
170  }
171  ret = av_opencl_buffer_write_image(deshake->opencl_ctx.cl_inbuf,
172  deshake->opencl_ctx.cl_inbuf_size,
173  0, in->data,deshake->opencl_ctx.in_plane_size,
174  deshake->opencl_ctx.plane_num);
175  if(ret < 0)
176  return ret;
177  return ret;
178 }
av_pix_fmt_desc_get
const AVPixFmtDescriptor * av_pix_fmt_desc_get(enum AVPixelFormat pix_fmt)
Definition: pixdesc.c:1778
AVFrame
This structure describes decoded (raw) audio or video data.
Definition: frame.h:76
in
About Git write you should know how to use GIT properly Luckily Git comes with excellent documentation git help man git shows you the available git< command > help man git< command > shows information about the subcommand< command > The most comprehensive manual is the website Git Reference visit they are quite exhaustive You do not need a special username or password All you need is to provide a ssh public key to the Git server admin What follows now is a basic introduction to Git and some FFmpeg specific guidelines Read it at least if you are granted commit privileges to the FFmpeg project you are expected to be familiar with these rules I if not You can get git from etc no matter how small Every one of them has been saved from looking like a fool by this many times It s very easy for stray debug output or cosmetic modifications to slip in
Definition: git-howto.txt:5
PLANE_NUM
#define PLANE_NUM
Definition: deshake_opencl.c:32
InterpolateMethod
InterpolateMethod
Definition: transform.h:39
AVFilterContext::inputs
AVFilterLink ** inputs
array of pointers to input links
Definition: avfilter.h:532
dict.h
Public dictionary API.
uint8_t
uint8_t
Definition: audio_convert.c:194
interpolate
static void interpolate(float *out, float v1, float v2, int size)
Definition: twinvq.c:280
ff_opencl_transform
int ff_opencl_transform(AVFilterContext *ctx, int width, int height, int cw, int ch, const float *matrix_y, const float *matrix_uv, enum InterpolateMethod interpolate, enum FillMethod fill, AVFrame *in, AVFrame *out)
Definition: deshake_opencl.c:48
AVPixFmtDescriptor::log2_chroma_h
uint8_t log2_chroma_h
Amount to shift the luma height right to find the chroma height.
Definition: pixdesc.h:75
AVFilterContext::priv
void * priv
private data for use by the filter
Definition: avfilter.h:545
av_log
void av_log(void *avcl, int level, const char *fmt,...)
Definition: log.c:246
link
Filter the word “frame” indicates either a video frame or a group of audio as stored in an AVFilterBuffer structure Format for each input and each output the list of supported formats For video that means pixel format For audio that means channel sample they are references to shared objects When the negotiation mechanism computes the intersection of the formats supported at each end of a link
Definition: filter_design.txt:23
FillMethod
FillMethod
Definition: transform.h:51
ret
ret
Definition: avfilter.c:821
av_opencl_buffer_create
int av_opencl_buffer_create(cl_mem *cl_buf, size_t cl_buf_size, int flags, void *host_ptr)
Create OpenCL buffer.
Definition: opencl.c:672
av_opencl_buffer_write_image
int av_opencl_buffer_write_image(cl_mem dst_cl_buf, size_t cl_buffer_size, int dst_cl_offset, uint8_t **src_data, int *plane_size, int plane_num)
Write image data from memory to OpenCL buffer.
Definition: opencl.c:742
ff_opencl_deshake_process_inout_buf
int ff_opencl_deshake_process_inout_buf(AVFilterContext *ctx, AVFrame *in, AVFrame *out)
Definition: deshake_opencl.c:136
NULL
NULL
Definition: eval.c:55
width
static int width
Definition: tests/utils.c:158
av_opencl_release_kernel
void av_opencl_release_kernel(AVOpenCLKernelEnv *env)
Release kernel object.
Definition: opencl.c:432
AVFrame::linesize
int linesize[AV_NUM_DATA_POINTERS]
For video, size in bytes of each picture line.
Definition: frame.h:101
TRANSFORM_OPENCL_SET_KERNEL_ARG
#define TRANSFORM_OPENCL_SET_KERNEL_ARG(arg_ptr)
Definition: deshake_opencl.c:41
AV_LOG_ERROR
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition: log.h:148
fill
spectrum bins to fill(1-based indexing) for m
av_opencl_create_kernel
int av_opencl_create_kernel(AVOpenCLKernelEnv *env, const char *kernel_name)
Create kernel object in the specified kernel environment.
Definition: opencl.c:390
height
BYTE int const BYTE int int int height
Definition: avisynth_c.h:713
av_opencl_buffer_release
void av_opencl_buffer_release(cl_mem *cl_buf)
Release OpenCL buffer.
Definition: opencl.c:683
deshake_opencl.h
AVERROR
Filter the word “frame” indicates either a video frame or a group of audio as stored in an AVFilterBuffer structure Format for each input and each output the list of supported formats For video that means pixel format For audio that means channel sample they are references to shared objects When the negotiation mechanism computes the intersection of the formats supported at each end of a all references to both lists are replaced with a reference to the intersection And when a single format is eventually chosen for a link amongst the remaining all references to the list are updated That means that if a filter requires that its input and output have the same format amongst a supported all it has to do is use a reference to the same list of formats query_formats can leave some formats unset and return AVERROR(EAGAIN) to cause the negotiation mechanism toagain later.That can be used by filters with complex requirements to use the format negotiated on one link to set the formats supported on another.Buffer references ownership and permissions
MATRIX_SIZE
#define MATRIX_SIZE
Definition: deshake_opencl.c:31
AVFrame::data
uint8_t * data[AV_NUM_DATA_POINTERS]
pointer to the picture/channel planes.
Definition: frame.h:87
ff_opencl_deshake_init
int ff_opencl_deshake_init(AVFilterContext *ctx)
Definition: deshake_opencl.c:97
common.h
common internal and external API header
av_opencl_buffer_write
int av_opencl_buffer_write(cl_mem dst_cl_buf, uint8_t *src_buf, size_t buf_size)
Write OpenCL buffer with data from src_buf.
Definition: opencl.c:696
av_opencl_uninit
void av_opencl_uninit(void)
Release OpenCL environment.
Definition: opencl.c:629
av_opencl_init
int av_opencl_init(AVOpenCLExternalEnv *ext_opencl_env)
Initialize the run time OpenCL environment and compile the kernel code registered with av_opencl_regi...
Definition: opencl.c:600
AVFilterContext
An instance of a filter.
Definition: avfilter.h:524
AVFrame::height
int height
Definition: frame.h:122
out
uint8_t pi<< 24) CONV_FUNC_GROUP(AV_SAMPLE_FMT_FLT, float, AV_SAMPLE_FMT_U8, uint8_t,(*(const uint8_t *) pi-0x80)*(1.0f/(1<< 7))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_DBL, double, AV_SAMPLE_FMT_U8, uint8_t,(*(const uint8_t *) pi-0x80)*(1.0/(1<< 7))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_U8, uint8_t, AV_SAMPLE_FMT_S16, int16_t,(*(const int16_t *) pi >> 8)+0x80) CONV_FUNC_GROUP(AV_SAMPLE_FMT_FLT, float, AV_SAMPLE_FMT_S16, int16_t,*(const int16_t *) pi *(1.0f/(1<< 15))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_DBL, double, AV_SAMPLE_FMT_S16, int16_t,*(const int16_t *) pi *(1.0/(1<< 15))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_U8, uint8_t, AV_SAMPLE_FMT_S32, int32_t,(*(const int32_t *) pi >> 24)+0x80) CONV_FUNC_GROUP(AV_SAMPLE_FMT_FLT, float, AV_SAMPLE_FMT_S32, int32_t,*(const int32_t *) pi *(1.0f/(1U<< 31))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_DBL, double, AV_SAMPLE_FMT_S32, int32_t,*(const int32_t *) pi *(1.0/(1U<< 31))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_U8, uint8_t, AV_SAMPLE_FMT_FLT, float, av_clip_uint8(lrintf(*(const float *) pi *(1<< 7))+0x80)) CONV_FUNC_GROUP(AV_SAMPLE_FMT_S16, int16_t, AV_SAMPLE_FMT_FLT, float, av_clip_int16(lrintf(*(const float *) pi *(1<< 15)))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_S32, int32_t, AV_SAMPLE_FMT_FLT, float, av_clipl_int32(llrintf(*(const float *) pi *(1U<< 31)))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_U8, uint8_t, AV_SAMPLE_FMT_DBL, double, av_clip_uint8(lrint(*(const double *) pi *(1<< 7))+0x80)) CONV_FUNC_GROUP(AV_SAMPLE_FMT_S16, int16_t, AV_SAMPLE_FMT_DBL, double, av_clip_int16(lrint(*(const double *) pi *(1<< 15)))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_S32, int32_t, AV_SAMPLE_FMT_DBL, double, av_clipl_int32(llrint(*(const double *) pi *(1U<< 31))))#define SET_CONV_FUNC_GROUP(ofmt, ifmt) static void set_generic_function(AudioConvert *ac){}void ff_audio_convert_free(AudioConvert **ac){if(!*ac) return;ff_dither_free(&(*ac) ->dc);av_freep(ac);}AudioConvert *ff_audio_convert_alloc(AVAudioResampleContext *avr, enum AVSampleFormat out_fmt, enum AVSampleFormat in_fmt, int channels, int sample_rate, int apply_map){AudioConvert *ac;int in_planar, out_planar;ac=av_mallocz(sizeof(*ac));if(!ac) return NULL;ac->avr=avr;ac->out_fmt=out_fmt;ac->in_fmt=in_fmt;ac->channels=channels;ac->apply_map=apply_map;if(avr->dither_method!=AV_RESAMPLE_DITHER_NONE &&av_get_packed_sample_fmt(out_fmt)==AV_SAMPLE_FMT_S16 &&av_get_bytes_per_sample(in_fmt) > 2){ac->dc=ff_dither_alloc(avr, out_fmt, in_fmt, channels, sample_rate, apply_map);if(!ac->dc){av_free(ac);return NULL;}return ac;}in_planar=av_sample_fmt_is_planar(in_fmt);out_planar=av_sample_fmt_is_planar(out_fmt);if(in_planar==out_planar){ac->func_type=CONV_FUNC_TYPE_FLAT;ac->planes=in_planar?ac->channels:1;}else if(in_planar) ac->func_type=CONV_FUNC_TYPE_INTERLEAVE;else ac->func_type=CONV_FUNC_TYPE_DEINTERLEAVE;set_generic_function(ac);if(ARCH_ARM) ff_audio_convert_init_arm(ac);if(ARCH_X86) ff_audio_convert_init_x86(ac);return ac;}int ff_audio_convert(AudioConvert *ac, AudioData *out, AudioData *in){int use_generic=1;int len=in->nb_samples;int p;if(ac->dc){av_dlog(ac->avr,"%d samples - audio_convert: %s to %s (dithered)\n", len, av_get_sample_fmt_name(ac->in_fmt), av_get_sample_fmt_name(ac->out_fmt));return ff_convert_dither(ac-> out
Definition: audio_convert.c:194
TRANSFORM_OPENCL_CHECK
#define TRANSFORM_OPENCL_CHECK(method,...)
Definition: deshake_opencl.c:34
ff_opencl_deshake_uninit
void ff_opencl_deshake_uninit(AVFilterContext *ctx)
Definition: deshake_opencl.c:124
av_opencl_buffer_read_image
int av_opencl_buffer_read_image(uint8_t **dst_data, int *plane_size, int plane_num, cl_mem src_cl_buf, size_t cl_buffer_size)
Read image data from OpenCL buffer.
Definition: opencl.c:783
Generated on Wed May 8 2024 06:52:03 for FFmpeg by   doxygen 1.8.11