lavu,lavfi,ffmpeg: Remove experimental OpenCL API

This was added in early 2013 and abandoned several months later; as far as
I can tell, there are no external users.  Future OpenCL use will be via
hwcontext, which requires neither special OpenCL-only API nor global state
in libavutil.

All internal users are also deleted - this is just the unsharp filter
(replaced by unsharp_opencl, which is more flexible) and the deshake filter
(no replacement).

14 files changed, 3 insertions, 1442 deletions

diff --git a/libavfilter/Makefile b/libavfilter/Makefile
index 30474ea77b..89737b5ad0 100644
--- a/libavfilter/Makefile
+++ b/libavfilter/Makefile
@@ -19,7 +19,6 @@ OBJS = allfilters.o                                                     \
        framequeue.o                                                     \
        graphdump.o                                                      \
        graphparser.o                                                    \
-       opencl_allkernels.o                                              \
        transform.o                                                      \
        video.o                                                          \
 
@@ -250,7 +249,6 @@ OBJS-$(CONFIG_NOISE_FILTER)                  += vf_noise.o
 OBJS-$(CONFIG_NULL_FILTER)                   += vf_null.o
 OBJS-$(CONFIG_OCR_FILTER)                    += vf_ocr.o
 OBJS-$(CONFIG_OCV_FILTER)                    += vf_libopencv.o
-OBJS-$(CONFIG_OPENCL)                        += deshake_opencl.o unsharp_opencl.o
 OBJS-$(CONFIG_OSCILLOSCOPE_FILTER)           += vf_datascope.o
 OBJS-$(CONFIG_OVERLAY_FILTER)                += vf_overlay.o framesync.o
 OBJS-$(CONFIG_OVERLAY_OPENCL_FILTER)         += vf_overlay_opencl.o opencl.o \
@@ -396,7 +394,6 @@ OBJS-$(CONFIG_MOVIE_FILTER)                  += src_movie.o
 SLIBOBJS-$(HAVE_GNU_WINDRES)                 += avfilterres.o
 
 SKIPHEADERS-$(CONFIG_LIBVIDSTAB)             += vidstabutils.h
-SKIPHEADERS-$(CONFIG_OPENCL)                 += opencl_internal.h deshake_opencl_kernel.h unsharp_opencl_kernel.h
 
 OBJS-$(CONFIG_SHARED)                        += log2_tab.o
 
diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c
index beacfce612..50e9d9a0cb 100644
--- a/libavfilter/allfilters.c
+++ b/libavfilter/allfilters.c
@@ -22,7 +22,6 @@
 #include "libavutil/thread.h"
 #include "avfilter.h"
 #include "config.h"
-#include "opencl_allkernels.h"
 
 
 #define REGISTER_FILTER(X, x, y)                                        \
@@ -407,7 +406,6 @@ static void register_all(void)
     REGISTER_FILTER_UNCONDITIONAL(vsink_buffer);
     REGISTER_FILTER_UNCONDITIONAL(af_afifo);
     REGISTER_FILTER_UNCONDITIONAL(vf_fifo);
-    ff_opencl_register_filter_kernel_code_all();
 }
 
 void avfilter_register_all(void)
diff --git a/libavfilter/deshake.h b/libavfilter/deshake.h
index 5e25bb3a7a..406cbab2f6 100644
--- a/libavfilter/deshake.h
+++ b/libavfilter/deshake.h
@@ -26,9 +26,6 @@
 #include "avfilter.h"
 #include "transform.h"
 #include "libavutil/pixelutils.h"
-#if CONFIG_OPENCL
-#include "libavutil/opencl.h"
-#endif
 
 
 enum SearchMethod {
@@ -53,24 +50,6 @@ typedef struct Transform {
     double zoom;          ///< Zoom percentage
 } Transform;
 
-#if CONFIG_OPENCL
-
-typedef struct DeshakeOpenclContext {
-    cl_command_queue command_queue;
-    cl_program program;
-    cl_kernel kernel_luma;
-    cl_kernel kernel_chroma;
-    int in_plane_size[8];
-    int out_plane_size[8];
-    int plane_num;
-    cl_mem cl_inbuf;
-    size_t cl_inbuf_size;
-    cl_mem cl_outbuf;
-    size_t cl_outbuf_size;
-} DeshakeOpenclContext;
-
-#endif
-
 #define MAX_R 64
 
 typedef struct DeshakeContext {
@@ -96,9 +75,6 @@ typedef struct DeshakeContext {
     int cy;
     char *filename;            ///< Motion search detailed log filename
     int opencl;
-#if CONFIG_OPENCL
-    DeshakeOpenclContext opencl_ctx;
-#endif
     int (* 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);
diff --git a/libavfilter/deshake_opencl.c b/libavfilter/deshake_opencl.c
deleted file mode 100644
index 877ec1dc03..0000000000
--- a/libavfilter/deshake_opencl.c
+++ /dev/null
@@ -1,198 +0,0 @@
-/*
- * Copyright (C) 2013 Wei Gao <weigao@multicorewareinc.com>
- * Copyright (C) 2013 Lenny Wang
- *
- * This file is part of FFmpeg.
- *
- * FFmpeg is free software; you can redistribute it and/or
- * modify it under the terms of the GNU Lesser General Public
- * License as published by the Free Software Foundation; either
- * version 2.1 of the License, or (at your option) any later version.
- *
- * FFmpeg is distributed in the hope that it will be useful,
- * but WITHOUT ANY WARRANTY; without even the implied warranty of
- * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
- * Lesser General Public License for more details.
- *
- * You should have received a copy of the GNU Lesser General Public
- * License along with FFmpeg; if not, write to the Free Software
- * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
- */
-
-/**
- * @file
- * transform input video
- */
-
-#include "libavutil/common.h"
-#include "libavutil/dict.h"
-#include "libavutil/pixdesc.h"
-#include "deshake_opencl.h"
-#include "libavutil/opencl_internal.h"
-
-#define PLANE_NUM 3
-#define ROUND_TO_16(a) (((((a) - 1)/16)+1)*16)
-
-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)
-{
-    int ret = 0;
-    cl_int status;
-    DeshakeContext *deshake = ctx->priv;
-    float4 packed_matrix_lu = {matrix_y[0], matrix_y[1], matrix_y[2], matrix_y[5]};
-    float4 packed_matrix_ch = {matrix_uv[0], matrix_uv[1], matrix_uv[2], matrix_uv[5]};
-    size_t global_worksize_lu[2] = {(size_t)ROUND_TO_16(width), (size_t)ROUND_TO_16(height)};
-    size_t global_worksize_ch[2] = {(size_t)ROUND_TO_16(cw), (size_t)(2*ROUND_TO_16(ch))};
-    size_t local_worksize[2] = {16, 16};
-    FFOpenclParam param_lu = {0};
-    FFOpenclParam param_ch = {0};
-    param_lu.ctx = param_ch.ctx = ctx;
-    param_lu.kernel = deshake->opencl_ctx.kernel_luma;
-    param_ch.kernel = deshake->opencl_ctx.kernel_chroma;
-
-    if ((unsigned int)interpolate > INTERPOLATE_BIQUADRATIC) {
-        av_log(ctx, AV_LOG_ERROR, "Selected interpolate method is invalid\n");
-        return AVERROR(EINVAL);
-    }
-    ret = avpriv_opencl_set_parameter(&param_lu,
-                                  FF_OPENCL_PARAM_INFO(deshake->opencl_ctx.cl_inbuf),
-                                  FF_OPENCL_PARAM_INFO(deshake->opencl_ctx.cl_outbuf),
-                                  FF_OPENCL_PARAM_INFO(packed_matrix_lu),
-                                  FF_OPENCL_PARAM_INFO(interpolate),
-                                  FF_OPENCL_PARAM_INFO(fill),
-                                  FF_OPENCL_PARAM_INFO(in->linesize[0]),
-                                  FF_OPENCL_PARAM_INFO(out->linesize[0]),
-                                  FF_OPENCL_PARAM_INFO(height),
-                                  FF_OPENCL_PARAM_INFO(width),
-                                  NULL);
-    if (ret < 0)
-        return ret;
-    ret = avpriv_opencl_set_parameter(&param_ch,
-                                  FF_OPENCL_PARAM_INFO(deshake->opencl_ctx.cl_inbuf),
-                                  FF_OPENCL_PARAM_INFO(deshake->opencl_ctx.cl_outbuf),
-                                  FF_OPENCL_PARAM_INFO(packed_matrix_ch),
-                                  FF_OPENCL_PARAM_INFO(interpolate),
-                                  FF_OPENCL_PARAM_INFO(fill),
-                                  FF_OPENCL_PARAM_INFO(in->linesize[0]),
-                                  FF_OPENCL_PARAM_INFO(out->linesize[0]),
-                                  FF_OPENCL_PARAM_INFO(in->linesize[1]),
-                                  FF_OPENCL_PARAM_INFO(out->linesize[1]),
-                                  FF_OPENCL_PARAM_INFO(height),
-                                  FF_OPENCL_PARAM_INFO(width),
-                                  FF_OPENCL_PARAM_INFO(ch),
-                                  FF_OPENCL_PARAM_INFO(cw),
-                                  NULL);
-    if (ret < 0)
-        return ret;
-    status = clEnqueueNDRangeKernel(deshake->opencl_ctx.command_queue,
-                                    deshake->opencl_ctx.kernel_luma, 2, NULL,
-                                    global_worksize_lu, local_worksize, 0, NULL, NULL);
-    status |= clEnqueueNDRangeKernel(deshake->opencl_ctx.command_queue,
-                                    deshake->opencl_ctx.kernel_chroma, 2, NULL,
-                                    global_worksize_ch, local_worksize, 0, NULL, NULL);
-    if (status != CL_SUCCESS) {
-        av_log(ctx, AV_LOG_ERROR, "OpenCL run kernel error occurred: %s\n", av_opencl_errstr(status));
-        return AVERROR_EXTERNAL;
-    }
-    ret = av_opencl_buffer_read_image(out->data, deshake->opencl_ctx.out_plane_size,
-                                      deshake->opencl_ctx.plane_num, deshake->opencl_ctx.cl_outbuf,
-                                      deshake->opencl_ctx.cl_outbuf_size);
-    if (ret < 0)
-        return ret;
-    return ret;
-}
-
-int ff_opencl_deshake_init(AVFilterContext *ctx)
-{
-    int ret = 0;
-    DeshakeContext *deshake = ctx->priv;
-    ret = av_opencl_init(NULL);
-    if (ret < 0)
-        return ret;
-    deshake->opencl_ctx.plane_num = PLANE_NUM;
-    deshake->opencl_ctx.command_queue = av_opencl_get_command_queue();
-    if (!deshake->opencl_ctx.command_queue) {
-        av_log(ctx, AV_LOG_ERROR, "Unable to get OpenCL command queue in filter 'deshake'\n");
-        return AVERROR(EINVAL);
-    }
-    deshake->opencl_ctx.program = av_opencl_compile("avfilter_transform", NULL);
-    if (!deshake->opencl_ctx.program) {
-        av_log(ctx, AV_LOG_ERROR, "OpenCL failed to compile program 'avfilter_transform'\n");
-        return AVERROR(EINVAL);
-    }
-    if (!deshake->opencl_ctx.kernel_luma) {
-        deshake->opencl_ctx.kernel_luma = clCreateKernel(deshake->opencl_ctx.program,
-                                                         "avfilter_transform_luma", &ret);
-        if (ret != CL_SUCCESS) {
-            av_log(ctx, AV_LOG_ERROR, "OpenCL failed to create kernel 'avfilter_transform_luma'\n");
-            return AVERROR(EINVAL);
-        }
-    }
-    if (!deshake->opencl_ctx.kernel_chroma) {
-        deshake->opencl_ctx.kernel_chroma = clCreateKernel(deshake->opencl_ctx.program,
-                                                           "avfilter_transform_chroma", &ret);
-        if (ret != CL_SUCCESS) {
-            av_log(ctx, AV_LOG_ERROR, "OpenCL failed to create kernel 'avfilter_transform_chroma'\n");
-            return AVERROR(EINVAL);
-        }
-    }
-    return ret;
-}
-
-void ff_opencl_deshake_uninit(AVFilterContext *ctx)
-{
-    DeshakeContext *deshake = ctx->priv;
-    av_opencl_buffer_release(&deshake->opencl_ctx.cl_inbuf);
-    av_opencl_buffer_release(&deshake->opencl_ctx.cl_outbuf);
-    clReleaseKernel(deshake->opencl_ctx.kernel_luma);
-    clReleaseKernel(deshake->opencl_ctx.kernel_chroma);
-    clReleaseProgram(deshake->opencl_ctx.program);
-    deshake->opencl_ctx.command_queue = NULL;
-    av_opencl_uninit();
-}
-
-int ff_opencl_deshake_process_inout_buf(AVFilterContext *ctx, AVFrame *in, AVFrame *out)
-{
-    int ret = 0;
-    AVFilterLink *link = ctx->inputs[0];
-    DeshakeContext *deshake = ctx->priv;
-    const int hshift = av_pix_fmt_desc_get(link->format)->log2_chroma_h;
-    int chroma_height = AV_CEIL_RSHIFT(link->h, hshift);
-
-    if ((!deshake->opencl_ctx.cl_inbuf) || (!deshake->opencl_ctx.cl_outbuf)) {
-        deshake->opencl_ctx.in_plane_size[0]  = (in->linesize[0] * in->height);
-        deshake->opencl_ctx.in_plane_size[1]  = (in->linesize[1] * chroma_height);
-        deshake->opencl_ctx.in_plane_size[2]  = (in->linesize[2] * chroma_height);
-        deshake->opencl_ctx.out_plane_size[0] = (out->linesize[0] * out->height);
-        deshake->opencl_ctx.out_plane_size[1] = (out->linesize[1] * chroma_height);
-        deshake->opencl_ctx.out_plane_size[2] = (out->linesize[2] * chroma_height);
-        deshake->opencl_ctx.cl_inbuf_size  = deshake->opencl_ctx.in_plane_size[0] +
-                                             deshake->opencl_ctx.in_plane_size[1] +
-                                             deshake->opencl_ctx.in_plane_size[2];
-        deshake->opencl_ctx.cl_outbuf_size = deshake->opencl_ctx.out_plane_size[0] +
-                                             deshake->opencl_ctx.out_plane_size[1] +
-                                             deshake->opencl_ctx.out_plane_size[2];
-        if (!deshake->opencl_ctx.cl_inbuf) {
-            ret = av_opencl_buffer_create(&deshake->opencl_ctx.cl_inbuf,
-                                            deshake->opencl_ctx.cl_inbuf_size,
-                                            CL_MEM_READ_ONLY, NULL);
-            if (ret < 0)
-                return ret;
-        }
-        if (!deshake->opencl_ctx.cl_outbuf) {
-            ret = av_opencl_buffer_create(&deshake->opencl_ctx.cl_outbuf,
-                                            deshake->opencl_ctx.cl_outbuf_size,
-                                            CL_MEM_READ_WRITE, NULL);
-            if (ret < 0)
-                return ret;
-        }
-    }
-    ret = av_opencl_buffer_write_image(deshake->opencl_ctx.cl_inbuf,
-                                 deshake->opencl_ctx.cl_inbuf_size,
-                                 0, in->data,deshake->opencl_ctx.in_plane_size,
-                                 deshake->opencl_ctx.plane_num);
-    return ret;
-}
diff --git a/libavfilter/deshake_opencl.h b/libavfilter/deshake_opencl.h
deleted file mode 100644
index f3d96dc4e9..0000000000
--- a/libavfilter/deshake_opencl.h
+++ /dev/null
@@ -1,45 +0,0 @@
-/*
- * Copyright (C) 2013 Wei Gao <weigao@multicorewareinc.com>
- *
- * This file is part of FFmpeg.
- *
- * FFmpeg is free software; you can redistribute it and/or
- * modify it under the terms of the GNU Lesser General Public
- * License as published by the Free Software Foundation; either
- * version 2.1 of the License, or (at your option) any later version.
- *
- * FFmpeg is distributed in the hope that it will be useful,
- * but WITHOUT ANY WARRANTY; without even the implied warranty of
- * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
- * Lesser General Public License for more details.
- *
- * You should have received a copy of the GNU Lesser General Public
- * License along with FFmpeg; if not, write to the Free Software
- * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
- */
-
-#ifndef AVFILTER_DESHAKE_OPENCL_H
-#define AVFILTER_DESHAKE_OPENCL_H
-
-#include "deshake.h"
-
-typedef struct float4 {
-    float x;
-    float y;
-    float z;
-    float w;
-} float4;
-
-int ff_opencl_deshake_init(AVFilterContext *ctx);
-
-void ff_opencl_deshake_uninit(AVFilterContext *ctx);
-
-int ff_opencl_deshake_process_inout_buf(AVFilterContext *ctx, AVFrame *in, AVFrame *out);
-
-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);
-
-#endif /* AVFILTER_DESHAKE_OPENCL_H */
diff --git a/libavfilter/deshake_opencl_kernel.h b/libavfilter/deshake_opencl_kernel.h
deleted file mode 100644
index dd45d6f60b..0000000000
--- a/libavfilter/deshake_opencl_kernel.h
+++ /dev/null
@@ -1,225 +0,0 @@
-/*
- * Copyright (C) 2013 Wei Gao <weigao@multicorewareinc.com>
- * Copyright (C) 2013 Lenny Wang
- *
- *
- * This file is part of FFmpeg.
- *
- * FFmpeg is free software; you can redistribute it and/or
- * modify it under the terms of the GNU Lesser General Public
- * License as published by the Free Software Foundation; either
- * version 2.1 of the License, or (at your option) any later version.
- *
- * FFmpeg is distributed in the hope that it will be useful,
- * but WITHOUT ANY WARRANTY; without even the implied warranty of
- * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
- * Lesser General Public License for more details.
- *
- * You should have received a copy of the GNU Lesser General Public
- * License along with FFmpeg; if not, write to the Free Software
- * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
- */
-
-#ifndef AVFILTER_DESHAKE_OPENCL_KERNEL_H
-#define AVFILTER_DESHAKE_OPENCL_KERNEL_H
-
-#include "libavutil/opencl.h"
-
-const char *ff_kernel_deshake_opencl = AV_OPENCL_KERNEL(
-inline unsigned char pixel(global const unsigned char *src, int x, int y,
-                           int w, int h,int stride, unsigned char def)
-{
-    return (x < 0 || y < 0 || x >= w || y >= h) ? def : src[x + y * stride];
-}
-
-unsigned char interpolate_nearest(float x, float y, global const unsigned char *src,
-                                  int width, int height, int stride, unsigned char def)
-{
-    return pixel(src, (int)(x + 0.5f), (int)(y + 0.5f), width, height, stride, def);
-}
-
-unsigned char interpolate_bilinear(float x, float y, global const unsigned char *src,
-                                   int width, int height, int stride, unsigned char def)
-{
-    int x_c, x_f, y_c, y_f;
-    int v1, v2, v3, v4;
-    x_f = (int)x;
-    y_f = (int)y;
-    x_c = x_f + 1;
-    y_c = y_f + 1;
-
-    if (x_f < -1 || x_f > width || y_f < -1 || y_f > height) {
-        return def;
-    } else {
-        v4 = pixel(src, x_f, y_f, width, height, stride, def);
-        v2 = pixel(src, x_c, y_f, width, height, stride, def);
-        v3 = pixel(src, x_f, y_c, width, height, stride, def);
-        v1 = pixel(src, x_c, y_c, width, height, stride, def);
-        return (v1*(x - x_f)*(y - y_f) + v2*((x - x_f)*(y_c - y)) +
-                v3*(x_c - x)*(y - y_f) + v4*((x_c - x)*(y_c - y)));
-    }
-}
-
-unsigned char interpolate_biquadratic(float x, float y, global const unsigned char *src,
-                                      int width, int height, int stride, unsigned char def)
-{
-    int     x_c, x_f, y_c, y_f;
-    unsigned char v1,  v2,  v3,  v4;
-    float   f1,  f2,  f3,  f4;
-    x_f = (int)x;
-    y_f = (int)y;
-    x_c = x_f + 1;
-    y_c = y_f + 1;
-
-    if (x_f < - 1 || x_f > width || y_f < -1 || y_f > height)
-        return def;
-    else {
-        v4 = pixel(src, x_f, y_f, width, height, stride, def);
-        v2 = pixel(src, x_c, y_f, width, height, stride, def);
-        v3 = pixel(src, x_f, y_c, width, height, stride, def);
-        v1 = pixel(src, x_c, y_c, width, height, stride, def);
-
-        f1 = 1 - sqrt((x_c - x) * (y_c - y));
-        f2 = 1 - sqrt((x_c - x) * (y - y_f));
-        f3 = 1 - sqrt((x - x_f) * (y_c - y));
-        f4 = 1 - sqrt((x - x_f) * (y - y_f));
-        return (v1 * f1 + v2 * f2 + v3 * f3 + v4 * f4) / (f1 + f2 + f3 + f4);
-    }
-}
-
-inline const float clipf(float a, float amin, float amax)
-{
-    if      (a < amin) return amin;
-    else if (a > amax) return amax;
-    else               return a;
-}
-
-inline int mirror(int v, int m)
-{
-    while ((unsigned)v > (unsigned)m) {
-        v = -v;
-        if (v < 0)
-            v += 2 * m;
-    }
-    return v;
-}
-
-kernel void avfilter_transform_luma(global unsigned char *src,
-                                    global unsigned char *dst,
-                                    float4 matrix,
-                                    int interpolate,
-                                    int fill,
-                                    int src_stride_lu,
-                                    int dst_stride_lu,
-                                    int height,
-                                    int width)
-{
-    int x = get_global_id(0);
-    int y = get_global_id(1);
-    int idx_dst = y * dst_stride_lu + x;
-    unsigned char def = 0;
-    float x_s = x * matrix.x + y * matrix.y + matrix.z;
-    float y_s = x * (-matrix.y) + y * matrix.x + matrix.w;
-
-    if (x < width && y < height) {
-        switch (fill) {
-            case 0: //FILL_BLANK
-                def = 0;
-                break;
-            case 1: //FILL_ORIGINAL
-                def = src[y*src_stride_lu + x];
-                break;
-            case 2: //FILL_CLAMP
-                y_s = clipf(y_s, 0, height - 1);
-                x_s = clipf(x_s, 0, width - 1);
-                def = src[(int)y_s * src_stride_lu + (int)x_s];
-                break;
-            case 3: //FILL_MIRROR
-                y_s = mirror(y_s, height - 1);
-                x_s = mirror(x_s, width - 1);
-                def = src[(int)y_s * src_stride_lu + (int)x_s];
-                break;
-        }
-        switch (interpolate) {
-            case 0: //INTERPOLATE_NEAREST
-                dst[idx_dst] = interpolate_nearest(x_s, y_s, src, width, height, src_stride_lu, def);
-                break;
-            case 1: //INTERPOLATE_BILINEAR
-                dst[idx_dst] = interpolate_bilinear(x_s, y_s, src, width, height, src_stride_lu, def);
-                break;
-            case 2: //INTERPOLATE_BIQUADRATIC
-                dst[idx_dst] = interpolate_biquadratic(x_s, y_s, src, width, height, src_stride_lu, def);
-                break;
-            default:
-                return;
-        }
-    }
-}
-
-kernel void avfilter_transform_chroma(global unsigned char *src,
-                                      global unsigned char *dst,
-                                      float4 matrix,
-                                      int interpolate,
-                                      int fill,
-                                      int src_stride_lu,
-                                      int dst_stride_lu,
-                                      int src_stride_ch,
-                                      int dst_stride_ch,
-                                      int height,
-                                      int width,
-                                      int ch,
-                                      int cw)
-{
-
-    int x = get_global_id(0);
-    int y = get_global_id(1);
-    int pad_ch = get_global_size(1)>>1;
-    global unsigned char *dst_u = dst + height * dst_stride_lu;
-    global unsigned char *src_u = src + height * src_stride_lu;
-    global unsigned char *dst_v = dst_u + ch * dst_stride_ch;
-    global unsigned char *src_v = src_u + ch * src_stride_ch;
-    src = y < pad_ch ? src_u : src_v;
-    dst = y < pad_ch ? dst_u : dst_v;
-    y = select(y - pad_ch, y, y < pad_ch);
-    float x_s = x * matrix.x + y * matrix.y + matrix.z;
-    float y_s = x * (-matrix.y) + y * matrix.x + matrix.w;
-    int idx_dst = y * dst_stride_ch + x;
-    unsigned char def;
-
-    if (x < cw && y < ch) {
-        switch (fill) {
-            case 0: //FILL_BLANK
-                def = 0;
-                break;
-            case 1: //FILL_ORIGINAL
-                def = src[y*src_stride_ch + x];
-                break;
-            case 2: //FILL_CLAMP
-                y_s = clipf(y_s, 0, ch - 1);
-                x_s = clipf(x_s, 0, cw - 1);
-                def = src[(int)y_s * src_stride_ch + (int)x_s];
-                break;
-            case 3: //FILL_MIRROR
-                y_s = mirror(y_s, ch - 1);
-                x_s = mirror(x_s, cw - 1);
-                def = src[(int)y_s * src_stride_ch + (int)x_s];
-                break;
-        }
-        switch (interpolate) {
-            case 0: //INTERPOLATE_NEAREST
-                dst[idx_dst] = interpolate_nearest(x_s, y_s, src, cw, ch, src_stride_ch, def);
-                break;
-            case 1: //INTERPOLATE_BILINEAR
-                dst[idx_dst] = interpolate_bilinear(x_s, y_s, src, cw, ch, src_stride_ch, def);
-                break;
-            case 2: //INTERPOLATE_BIQUADRATIC
-                dst[idx_dst] = interpolate_biquadratic(x_s, y_s, src, cw, ch, src_stride_ch, def);
-                break;
-            default:
-                return;
-        }
-    }
-}
-);
-
-#endif /* AVFILTER_DESHAKE_OPENCL_KERNEL_H */
diff --git a/libavfilter/opencl_allkernels.c b/libavfilter/opencl_allkernels.c
deleted file mode 100644
index 6d80fa8598..0000000000
--- a/libavfilter/opencl_allkernels.c
+++ /dev/null
@@ -1,41 +0,0 @@
-/*
- * Copyright (C) 2013 Wei Gao <weigao@multicorewareinc.com>
- *
- * This file is part of FFmpeg.
- *
- * FFmpeg is free software; you can redistribute it and/or
- * modify it under the terms of the GNU Lesser General Public
- * License as published by the Free Software Foundation; either
- * version 2.1 of the License, or (at your option) any later version.
- *
- * FFmpeg is distributed in the hope that it will be useful,
- * but WITHOUT ANY WARRANTY; without even the implied warranty of
- * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
- * Lesser General Public License for more details.
- *
- * You should have received a copy of the GNU Lesser General Public
- * License along with FFmpeg; if not, write to the Free Software
- * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
- */
-
-#include "opencl_allkernels.h"
-#if CONFIG_OPENCL
-#include "libavutil/opencl.h"
-#include "deshake_opencl_kernel.h"
-#include "unsharp_opencl_kernel.h"
-#endif
-
-#define OPENCL_REGISTER_KERNEL_CODE(X, x)                                              \
-    {                                                                                  \
-        if (CONFIG_##X##_FILTER) {                                                     \
-            av_opencl_register_kernel_code(ff_kernel_##x##_opencl);                    \
-        }                                                                              \
-    }
-
-void ff_opencl_register_filter_kernel_code_all(void)
-{
- #if CONFIG_OPENCL
-   OPENCL_REGISTER_KERNEL_CODE(DESHAKE,     deshake);
-   OPENCL_REGISTER_KERNEL_CODE(UNSHARP,     unsharp);
- #endif
-}
diff --git a/libavfilter/opencl_allkernels.h b/libavfilter/opencl_allkernels.h
deleted file mode 100644
index 57b650d27b..0000000000
--- a/libavfilter/opencl_allkernels.h
+++ /dev/null
@@ -1,29 +0,0 @@
-/*
- * Copyright (C) 2013 Wei Gao <weigao@multicorewareinc.com>
- *
- * This file is part of FFmpeg.
- *
- * FFmpeg is free software; you can redistribute it and/or
- * modify it under the terms of the GNU Lesser General Public
- * License as published by the Free Software Foundation; either
- * version 2.1 of the License, or (at your option) any later version.
- *
- * FFmpeg is distributed in the hope that it will be useful,
- * but WITHOUT ANY WARRANTY; without even the implied warranty of
- * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
- * Lesser General Public License for more details.
- *
- * You should have received a copy of the GNU Lesser General Public
- * License along with FFmpeg; if not, write to the Free Software
- * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
- */
-
-#ifndef AVFILTER_OPENCL_ALLKERNELS_H
-#define AVFILTER_OPENCL_ALLKERNELS_H
-
-#include "avfilter.h"
-#include "config.h"
-
-void ff_opencl_register_filter_kernel_code_all(void);
-
-#endif /* AVFILTER_OPENCL_ALLKERNELS_H */
diff --git a/libavfilter/unsharp.h b/libavfilter/unsharp.h
index 340a6a0caf..caff986fc1 100644
--- a/libavfilter/unsharp.h
+++ b/libavfilter/unsharp.h
@@ -24,38 +24,10 @@
 
 #include "config.h"
 #include "avfilter.h"
-#if CONFIG_OPENCL
-#include "libavutil/opencl.h"
-#endif
 
 #define MIN_MATRIX_SIZE 3
 #define MAX_MATRIX_SIZE 63
 
-#if CONFIG_OPENCL
-
-typedef struct UnsharpOpenclContext {
-    cl_command_queue command_queue;
-    cl_program program;
-    cl_kernel kernel_default;
-    cl_kernel kernel_luma;
-    cl_kernel kernel_chroma;
-    cl_mem cl_luma_mask;
-    cl_mem cl_chroma_mask;
-    cl_mem cl_luma_mask_x;
-    cl_mem cl_chroma_mask_x;
-    cl_mem cl_luma_mask_y;
-    cl_mem cl_chroma_mask_y;
-    int in_plane_size[8];
-    int out_plane_size[8];
-    int plane_num;
-    cl_mem cl_inbuf;
-    size_t cl_inbuf_size;
-    cl_mem cl_outbuf;
-    size_t cl_outbuf_size;
-    int use_fast_kernels;
-} UnsharpOpenclContext;
-
-#endif
 
 typedef struct UnsharpFilterParam {
     int msize_x;                             ///< matrix width
@@ -76,9 +48,6 @@ typedef struct UnsharpContext {
     UnsharpFilterParam chroma; ///< chroma parameters (width, height, amount)
     int hsub, vsub;
     int opencl;
-#if CONFIG_OPENCL
-    UnsharpOpenclContext opencl_ctx;
-#endif
     int (* apply_unsharp)(AVFilterContext *ctx, AVFrame *in, AVFrame *out);
 } UnsharpContext;
 
diff --git a/libavfilter/unsharp_opencl.c b/libavfilter/unsharp_opencl.c
deleted file mode 100644
index 1545455846..0000000000
--- a/libavfilter/unsharp_opencl.c
+++ /dev/null
@@ -1,422 +0,0 @@
-/*
- * Copyright (C) 2013 Wei Gao <weigao@multicorewareinc.com>
- * Copyright (C) 2013 Lenny Wang
- *
- * This file is part of FFmpeg.
- *
- * FFmpeg is free software; you can redistribute it and/or
- * modify it under the terms of the GNU Lesser General Public
- * License as published by the Free Software Foundation; either
- * version 2.1 of the License, or (at your option) any later version.
- *
- * FFmpeg is distributed in the hope that it will be useful,
- * but WITHOUT ANY WARRANTY; without even the implied warranty of
- * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
- * Lesser General Public License for more details.
- *
- * You should have received a copy of the GNU Lesser General Public
- * License along with FFmpeg; if not, write to the Free Software
- * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
- */
-
-/**
- * @file
- * unsharp input video
- */
-
-#include "unsharp_opencl.h"
-#include "libavutil/common.h"
-#include "libavutil/opencl_internal.h"
-
-#define PLANE_NUM 3
-#define ROUND_TO_16(a) (((((a) - 1)/16)+1)*16)
-
-static inline void add_mask_counter(uint32_t *dst, uint32_t *counter1, uint32_t *counter2, int len)
-{
-    int i;
-    for (i = 0; i < len; i++) {
-        dst[i] = counter1[i] + counter2[i];
-    }
-}
-
-static int compute_mask(int step, uint32_t *mask)
-{
-    int i, z, ret = 0;
-    int counter_size = sizeof(uint32_t) * (2 * step + 1);
-    uint32_t *temp1_counter, *temp2_counter, **counter = NULL;
-    temp1_counter = av_mallocz(counter_size);
-    if (!temp1_counter) {
-        ret = AVERROR(ENOMEM);
-        goto end;
-    }
-    temp2_counter = av_mallocz(counter_size);
-    if (!temp2_counter) {
-        ret = AVERROR(ENOMEM);
-        goto end;
-    }
-    counter = av_mallocz_array(2 * step + 1, sizeof(uint32_t *));
-    if (!counter) {
-        ret = AVERROR(ENOMEM);
-        goto end;
-    }
-    for (i = 0; i < 2 * step + 1; i++) {
-        counter[i] = av_mallocz(counter_size);
-        if (!counter[i]) {
-            ret = AVERROR(ENOMEM);
-            goto end;
-        }
-    }
-    for (i = 0; i < 2 * step + 1; i++) {
-        memset(temp1_counter, 0, counter_size);
-        temp1_counter[i] = 1;
-        for (z = 0; z < step * 2; z += 2) {
-            add_mask_counter(temp2_counter, counter[z], temp1_counter, step * 2);
-            memcpy(counter[z], temp1_counter, counter_size);
-            add_mask_counter(temp1_counter, counter[z + 1], temp2_counter, step * 2);
-            memcpy(counter[z + 1], temp2_counter, counter_size);
-        }
-    }
-    memcpy(mask, temp1_counter, counter_size);
-end:
-    av_freep(&temp1_counter);
-    av_freep(&temp2_counter);
-    for (i = 0; counter && i < 2 * step + 1; i++) {
-        av_freep(&counter[i]);
-    }
-    av_freep(&counter);
-    return ret;
-}
-
-static int copy_separable_masks(cl_mem cl_mask_x, cl_mem cl_mask_y, int step_x, int step_y)
-{
-    int ret = 0;
-    uint32_t *mask_x, *mask_y;
-    size_t size_mask_x = sizeof(uint32_t) * (2 * step_x + 1);
-    size_t size_mask_y = sizeof(uint32_t) * (2 * step_y + 1);
-    mask_x = av_mallocz_array(2 * step_x + 1, sizeof(uint32_t));
-    if (!mask_x) {
-        ret = AVERROR(ENOMEM);
-        goto end;
-    }
-    mask_y = av_mallocz_array(2 * step_y + 1, sizeof(uint32_t));
-    if (!mask_y) {
-        ret = AVERROR(ENOMEM);
-        goto end;
-    }
-
-    ret = compute_mask(step_x, mask_x);
-    if (ret < 0)
-        goto end;
-    ret = compute_mask(step_y, mask_y);
-    if (ret < 0)
-        goto end;
-
-    ret = av_opencl_buffer_write(cl_mask_x, (uint8_t *)mask_x, size_mask_x);
-    ret = av_opencl_buffer_write(cl_mask_y, (uint8_t *)mask_y, size_mask_y);
-end:
-    av_freep(&mask_x);
-    av_freep(&mask_y);
-
-    return ret;
-}
-
-static int generate_mask(AVFilterContext *ctx)
-{
-    cl_mem masks[4];
-    cl_mem mask_matrix[2];
-    int i, ret = 0, step_x[2], step_y[2];
-
-    UnsharpContext *unsharp = ctx->priv;
-    mask_matrix[0] = unsharp->opencl_ctx.cl_luma_mask;
-    mask_matrix[1] = unsharp->opencl_ctx.cl_chroma_mask;
-    masks[0] = unsharp->opencl_ctx.cl_luma_mask_x;
-    masks[1] = unsharp->opencl_ctx.cl_luma_mask_y;
-    masks[2] = unsharp->opencl_ctx.cl_chroma_mask_x;
-    masks[3] = unsharp->opencl_ctx.cl_chroma_mask_y;
-    step_x[0] = unsharp->luma.steps_x;
-    step_x[1] = unsharp->chroma.steps_x;
-    step_y[0] = unsharp->luma.steps_y;
-    step_y[1] = unsharp->chroma.steps_y;
-
-    /* use default kernel if any matrix dim larger than 8 due to limited local mem size */
-    if (step_x[0]>8 || step_x[1]>8 || step_y[0]>8 || step_y[1]>8)
-        unsharp->opencl_ctx.use_fast_kernels = 0;
-    else
-        unsharp->opencl_ctx.use_fast_kernels = 1;
-
-    if (!masks[0] || !masks[1] || !masks[2] || !masks[3]) {
-        av_log(ctx, AV_LOG_ERROR, "Luma mask and chroma mask should not be NULL\n");
-        return AVERROR(EINVAL);
-    }
-    if (!mask_matrix[0] || !mask_matrix[1]) {
-        av_log(ctx, AV_LOG_ERROR, "Luma mask and chroma mask should not be NULL\n");
-        return AVERROR(EINVAL);
-    }
-    for (i = 0; i < 2; i++) {
-        ret = copy_separable_masks(masks[2*i], masks[2*i+1], step_x[i], step_y[i]);
-        if (ret < 0)
-            return ret;
-    }
-    return ret;
-}
-
-int ff_opencl_apply_unsharp(AVFilterContext *ctx, AVFrame *in, AVFrame *out)
-{
-    int ret;
-    AVFilterLink *link = ctx->inputs[0];
-    UnsharpContext *unsharp = ctx->priv;
-    cl_int status;
-    FFOpenclParam kernel1 = {0};
-    FFOpenclParam kernel2 = {0};
-    int width = link->w;
-    int height = link->h;
-    int cw = AV_CEIL_RSHIFT(link->w, unsharp->hsub);
-    int ch = AV_CEIL_RSHIFT(link->h, unsharp->vsub);
-    size_t globalWorkSize1d = width * height + 2 * ch * cw;
-    size_t globalWorkSize2dLuma[2];
-    size_t globalWorkSize2dChroma[2];
-    size_t localWorkSize2d[2] = {16, 16};
-
-    if (unsharp->opencl_ctx.use_fast_kernels) {
-        globalWorkSize2dLuma[0] = (size_t)ROUND_TO_16(width);
-        globalWorkSize2dLuma[1] = (size_t)ROUND_TO_16(height);
-        globalWorkSize2dChroma[0] = (size_t)ROUND_TO_16(cw);
-        globalWorkSize2dChroma[1] = (size_t)(2*ROUND_TO_16(ch));
-
-        kernel1.ctx = ctx;
-        kernel1.kernel = unsharp->opencl_ctx.kernel_luma;
-        ret = avpriv_opencl_set_parameter(&kernel1,
-                                      FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_inbuf),
-                                      FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_outbuf),
-                                      FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_luma_mask_x),
-                                      FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_luma_mask_y),
-                                      FF_OPENCL_PARAM_INFO(unsharp->luma.amount),
-                                      FF_OPENCL_PARAM_INFO(unsharp->luma.scalebits),
-                                      FF_OPENCL_PARAM_INFO(unsharp->luma.halfscale),
-                                      FF_OPENCL_PARAM_INFO(in->linesize[0]),
-                                      FF_OPENCL_PARAM_INFO(out->linesize[0]),
-                                      FF_OPENCL_PARAM_INFO(width),
-                                      FF_OPENCL_PARAM_INFO(height),
-                                      NULL);
-        if (ret < 0)
-            return ret;
-
-        kernel2.ctx = ctx;
-        kernel2.kernel = unsharp->opencl_ctx.kernel_chroma;
-        ret = avpriv_opencl_set_parameter(&kernel2,
-                                      FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_inbuf),
-                                      FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_outbuf),
-                                      FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_chroma_mask_x),
-                                      FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_chroma_mask_y),
-                                      FF_OPENCL_PARAM_INFO(unsharp->chroma.amount),
-                                      FF_OPENCL_PARAM_INFO(unsharp->chroma.scalebits),
-                                      FF_OPENCL_PARAM_INFO(unsharp->chroma.halfscale),
-                                      FF_OPENCL_PARAM_INFO(in->linesize[0]),
-                                      FF_OPENCL_PARAM_INFO(in->linesize[1]),
-                                      FF_OPENCL_PARAM_INFO(out->linesize[0]),
-                                      FF_OPENCL_PARAM_INFO(out->linesize[1]),
-                                      FF_OPENCL_PARAM_INFO(link->w),
-                                      FF_OPENCL_PARAM_INFO(link->h),
-                                      FF_OPENCL_PARAM_INFO(cw),
-                                      FF_OPENCL_PARAM_INFO(ch),
-                                      NULL);
-        if (ret < 0)
-            return ret;
-        status = clEnqueueNDRangeKernel(unsharp->opencl_ctx.command_queue,
-                                        unsharp->opencl_ctx.kernel_luma, 2, NULL,
-                                        globalWorkSize2dLuma, localWorkSize2d, 0, NULL, NULL);
-        status |=clEnqueueNDRangeKernel(unsharp->opencl_ctx.command_queue,
-                                        unsharp->opencl_ctx.kernel_chroma, 2, NULL,
-                                        globalWorkSize2dChroma, localWorkSize2d, 0, NULL, NULL);
-        if (status != CL_SUCCESS) {
-            av_log(ctx, AV_LOG_ERROR, "OpenCL run kernel error occurred: %s\n", av_opencl_errstr(status));
-            return AVERROR_EXTERNAL;
-        }
-    } else {    /* use default kernel */
-        kernel1.ctx = ctx;
-        kernel1.kernel = unsharp->opencl_ctx.kernel_default;
-
-        ret = avpriv_opencl_set_parameter(&kernel1,
-                                      FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_inbuf),
-                                      FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_outbuf),
-                                      FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_luma_mask),
-                                      FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_chroma_mask),
-                                      FF_OPENCL_PARAM_INFO(unsharp->luma.amount),
-                                      FF_OPENCL_PARAM_INFO(unsharp->chroma.amount),
-                                      FF_OPENCL_PARAM_INFO(unsharp->luma.steps_x),
-                                      FF_OPENCL_PARAM_INFO(unsharp->luma.steps_y),
-                                      FF_OPENCL_PARAM_INFO(unsharp->chroma.steps_x),
-                                      FF_OPENCL_PARAM_INFO(unsharp->chroma.steps_y),
-                                      FF_OPENCL_PARAM_INFO(unsharp->luma.scalebits),
-                                      FF_OPENCL_PARAM_INFO(unsharp->chroma.scalebits),
-                                      FF_OPENCL_PARAM_INFO(unsharp->luma.halfscale),
-                                      FF_OPENCL_PARAM_INFO(unsharp->chroma.halfscale),
-                                      FF_OPENCL_PARAM_INFO(in->linesize[0]),
-                                      FF_OPENCL_PARAM_INFO(in->linesize[1]),
-                                      FF_OPENCL_PARAM_INFO(out->linesize[0]),
-                                      FF_OPENCL_PARAM_INFO(out->linesize[1]),
-                                      FF_OPENCL_PARAM_INFO(link->h),
-                                      FF_OPENCL_PARAM_INFO(link->w),
-                                      FF_OPENCL_PARAM_INFO(ch),
-                                      FF_OPENCL_PARAM_INFO(cw),
-                                      NULL);
-        if (ret < 0)
-            return ret;
-        status = clEnqueueNDRangeKernel(unsharp->opencl_ctx.command_queue,
-                                        unsharp->opencl_ctx.kernel_default, 1, NULL,
-                                        &globalWorkSize1d, NULL, 0, NULL, NULL);
-        if (status != CL_SUCCESS) {
-            av_log(ctx, AV_LOG_ERROR, "OpenCL run kernel error occurred: %s\n", av_opencl_errstr(status));
-            return AVERROR_EXTERNAL;
-        }
-    }
-    //blocking map is suffficient, no need for clFinish
-    //clFinish(unsharp->opencl_ctx.command_queue);
-
-    return av_opencl_buffer_read_image(out->data, unsharp->opencl_ctx.out_plane_size,
-                                       unsharp->opencl_ctx.plane_num, unsharp->opencl_ctx.cl_outbuf,
-                                       unsharp->opencl_ctx.cl_outbuf_size);
-}
-
-int ff_opencl_unsharp_init(AVFilterContext *ctx)
-{
-    int ret = 0;
-    char build_opts[96];
-    UnsharpContext *unsharp = ctx->priv;
-    ret = av_opencl_init(NULL);
-    if (ret < 0)
-        return ret;
-    ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_luma_mask,
-                                  sizeof(uint32_t) * (2 * unsharp->luma.steps_x + 1) * (2 * unsharp->luma.steps_y + 1),
-                                  CL_MEM_READ_ONLY, NULL);
-    if (ret < 0)
-        return ret;
-    ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_chroma_mask,
-                                  sizeof(uint32_t) * (2 * unsharp->chroma.steps_x + 1) * (2 * unsharp->chroma.steps_y + 1),
-                                  CL_MEM_READ_ONLY, NULL);
-    // separable filters
-    if (ret < 0)
-        return ret;
-    ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_luma_mask_x,
-                                  sizeof(uint32_t) * (2 * unsharp->luma.steps_x + 1),
-                                  CL_MEM_READ_ONLY, NULL);
-    if (ret < 0)
-        return ret;
-    ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_luma_mask_y,
-                                  sizeof(uint32_t) * (2 * unsharp->luma.steps_y + 1),
-                                  CL_MEM_READ_ONLY, NULL);
-    if (ret < 0)
-        return ret;
-    ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_chroma_mask_x,
-                                  sizeof(uint32_t) * (2 * unsharp->chroma.steps_x + 1),
-                                  CL_MEM_READ_ONLY, NULL);
-    if (ret < 0)
-        return ret;
-    ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_chroma_mask_y,
-                                  sizeof(uint32_t) * (2 * unsharp->chroma.steps_y + 1),
-                                  CL_MEM_READ_ONLY, NULL);
-    if (ret < 0)
-        return ret;
-    ret = generate_mask(ctx);
-    if (ret < 0)
-        return ret;
-    unsharp->opencl_ctx.plane_num = PLANE_NUM;
-    unsharp->opencl_ctx.command_queue = av_opencl_get_command_queue();
-    if (!unsharp->opencl_ctx.command_queue) {
-        av_log(ctx, AV_LOG_ERROR, "Unable to get OpenCL command queue in filter 'unsharp'\n");
-        return AVERROR(EINVAL);
-    }
-    snprintf(build_opts, 96, "-D LU_RADIUS_X=%d -D LU_RADIUS_Y=%d -D CH_RADIUS_X=%d -D CH_RADIUS_Y=%d",
-            2*unsharp->luma.steps_x+1, 2*unsharp->luma.steps_y+1, 2*unsharp->chroma.steps_x+1, 2*unsharp->chroma.steps_y+1);
-    unsharp->opencl_ctx.program = av_opencl_compile("unsharp", build_opts);
-    if (!unsharp->opencl_ctx.program) {
-        av_log(ctx, AV_LOG_ERROR, "OpenCL failed to compile program 'unsharp'\n");
-        return AVERROR(EINVAL);
-    }
-    if (unsharp->opencl_ctx.use_fast_kernels) {
-        if (!unsharp->opencl_ctx.kernel_luma) {
-            unsharp->opencl_ctx.kernel_luma = clCreateKernel(unsharp->opencl_ctx.program, "unsharp_luma", &ret);
-            if (ret != CL_SUCCESS) {
-                av_log(ctx, AV_LOG_ERROR, "OpenCL failed to create kernel 'unsharp_luma'\n");
-                return ret;
-            }
-        }
-        if (!unsharp->opencl_ctx.kernel_chroma) {
-            unsharp->opencl_ctx.kernel_chroma = clCreateKernel(unsharp->opencl_ctx.program, "unsharp_chroma", &ret);
-            if (ret < 0) {
-                av_log(ctx, AV_LOG_ERROR, "OpenCL failed to create kernel 'unsharp_chroma'\n");
-                return ret;
-            }
-        }
-    }
-    else {
-        if (!unsharp->opencl_ctx.kernel_default) {
-            unsharp->opencl_ctx.kernel_default = clCreateKernel(unsharp->opencl_ctx.program, "unsharp_default", &ret);
-            if (ret < 0) {
-                av_log(ctx, AV_LOG_ERROR, "OpenCL failed to create kernel 'unsharp_default'\n");
-                return ret;
-            }
-        }
-    }
-    return ret;
-}
-
-void ff_opencl_unsharp_uninit(AVFilterContext *ctx)
-{
-    UnsharpContext *unsharp = ctx->priv;
-    av_opencl_buffer_release(&unsharp->opencl_ctx.cl_inbuf);
-    av_opencl_buffer_release(&unsharp->opencl_ctx.cl_outbuf);
-    av_opencl_buffer_release(&unsharp->opencl_ctx.cl_luma_mask);
-    av_opencl_buffer_release(&unsharp->opencl_ctx.cl_chroma_mask);
-    av_opencl_buffer_release(&unsharp->opencl_ctx.cl_luma_mask_x);
-    av_opencl_buffer_release(&unsharp->opencl_ctx.cl_chroma_mask_x);
-    av_opencl_buffer_release(&unsharp->opencl_ctx.cl_luma_mask_y);
-    av_opencl_buffer_release(&unsharp->opencl_ctx.cl_chroma_mask_y);
-    clReleaseKernel(unsharp->opencl_ctx.kernel_default);
-    clReleaseKernel(unsharp->opencl_ctx.kernel_luma);
-    clReleaseKernel(unsharp->opencl_ctx.kernel_chroma);
-    clReleaseProgram(unsharp->opencl_ctx.program);
-    unsharp->opencl_ctx.command_queue = NULL;
-    av_opencl_uninit();
-}
-
-int ff_opencl_unsharp_process_inout_buf(AVFilterContext *ctx, AVFrame *in, AVFrame *out)
-{
-    int ret = 0;
-    AVFilterLink *link = ctx->inputs[0];
-    UnsharpContext *unsharp = ctx->priv;
-    int ch = AV_CEIL_RSHIFT(link->h, unsharp->vsub);
-
-    if ((!unsharp->opencl_ctx.cl_inbuf) || (!unsharp->opencl_ctx.cl_outbuf)) {
-        unsharp->opencl_ctx.in_plane_size[0]  = (in->linesize[0] * in->height);
-        unsharp->opencl_ctx.in_plane_size[1]  = (in->linesize[1] * ch);
-        unsharp->opencl_ctx.in_plane_size[2]  = (in->linesize[2] * ch);
-        unsharp->opencl_ctx.out_plane_size[0] = (out->linesize[0] * out->height);
-        unsharp->opencl_ctx.out_plane_size[1] = (out->linesize[1] * ch);
-        unsharp->opencl_ctx.out_plane_size[2] = (out->linesize[2] * ch);
-        unsharp->opencl_ctx.cl_inbuf_size  = unsharp->opencl_ctx.in_plane_size[0] +
-                                             unsharp->opencl_ctx.in_plane_size[1] +
-                                             unsharp->opencl_ctx.in_plane_size[2];
-        unsharp->opencl_ctx.cl_outbuf_size = unsharp->opencl_ctx.out_plane_size[0] +
-                                             unsharp->opencl_ctx.out_plane_size[1] +
-                                             unsharp->opencl_ctx.out_plane_size[2];
-        if (!unsharp->opencl_ctx.cl_inbuf) {
-            ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_inbuf,
-                                          unsharp->opencl_ctx.cl_inbuf_size,
-                                          CL_MEM_READ_ONLY, NULL);
-            if (ret < 0)
-                return ret;
-        }
-        if (!unsharp->opencl_ctx.cl_outbuf) {
-            ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_outbuf,
-                                          unsharp->opencl_ctx.cl_outbuf_size,
-                                          CL_MEM_READ_WRITE, NULL);
-            if (ret < 0)
-                return ret;
-        }
-    }
-    return av_opencl_buffer_write_image(unsharp->opencl_ctx.cl_inbuf,
-                                        unsharp->opencl_ctx.cl_inbuf_size,
-                                        0, in->data, unsharp->opencl_ctx.in_plane_size,
-                                        unsharp->opencl_ctx.plane_num);
-}
diff --git a/libavfilter/unsharp_opencl.h b/libavfilter/unsharp_opencl.h
deleted file mode 100644
index 3aefab62e0..0000000000
--- a/libavfilter/unsharp_opencl.h
+++ /dev/null
@@ -1,34 +0,0 @@
-/*
- * Copyright (C) 2013 Wei Gao <weigao@multicorewareinc.com>
- *
- * This file is part of FFmpeg.
- *
- * FFmpeg is free software; you can redistribute it and/or
- * modify it under the terms of the GNU Lesser General Public
- * License as published by the Free Software Foundation; either
- * version 2.1 of the License, or (at your option) any later version.
- *
- * FFmpeg is distributed in the hope that it will be useful,
- * but WITHOUT ANY WARRANTY; without even the implied warranty of
- * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
- * Lesser General Public License for more details.
- *
- * You should have received a copy of the GNU Lesser General Public
- * License along with FFmpeg; if not, write to the Free Software
- * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
- */
-
-#ifndef AVFILTER_UNSHARP_OPENCL_H
-#define AVFILTER_UNSHARP_OPENCL_H
-
-#include "unsharp.h"
-
-int ff_opencl_unsharp_init(AVFilterContext *ctx);
-
-void ff_opencl_unsharp_uninit(AVFilterContext *ctx);
-
-int ff_opencl_unsharp_process_inout_buf(AVFilterContext *ctx, AVFrame *in, AVFrame *out);
-
-int ff_opencl_apply_unsharp(AVFilterContext *ctx, AVFrame *in, AVFrame *out);
-
-#endif /* AVFILTER_UNSHARP_OPENCL_H */
diff --git a/libavfilter/unsharp_opencl_kernel.h b/libavfilter/unsharp_opencl_kernel.h
deleted file mode 100644
index 307d0f1814..0000000000
--- a/libavfilter/unsharp_opencl_kernel.h
+++ /dev/null
@@ -1,342 +0,0 @@
-/*
- * Copyright (C) 2013 Wei Gao <weigao@multicorewareinc.com>
- * Copyright (C) 2013 Lenny Wang
- *
- * This file is part of FFmpeg.
- *
- * FFmpeg is free software; you can redistribute it and/or
- * modify it under the terms of the GNU Lesser General Public
- * License as published by the Free Software Foundation; either
- * version 2.1 of the License, or (at your option) any later version.
- *
- * FFmpeg is distributed in the hope that it will be useful,
- * but WITHOUT ANY WARRANTY; without even the implied warranty of
- * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
- * Lesser General Public License for more details.
- *
- * You should have received a copy of the GNU Lesser General Public
- * License along with FFmpeg; if not, write to the Free Software
- * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
- */
-
-#ifndef AVFILTER_UNSHARP_OPENCL_KERNEL_H
-#define AVFILTER_UNSHARP_OPENCL_KERNEL_H
-
-#include "libavutil/opencl.h"
-
-const char *ff_kernel_unsharp_opencl = AV_OPENCL_KERNEL(
-inline unsigned char clip_uint8(int a)
-{
-    if (a & (~0xFF))
-        return (-a)>>31;
-    else
-        return a;
-}
-
-kernel void unsharp_luma(
-                    global unsigned char *src,
-                    global unsigned char *dst,
-                    global int *mask_x,
-                    global int *mask_y,
-                    int amount,
-                    int scalebits,
-                    int halfscale,
-                    int src_stride,
-                    int dst_stride,
-                    int width,
-                    int height)
-{
-    int2 threadIdx, blockIdx, globalIdx;
-    threadIdx.x = get_local_id(0);
-    threadIdx.y = get_local_id(1);
-    blockIdx.x = get_group_id(0);
-    blockIdx.y = get_group_id(1);
-    globalIdx.x = get_global_id(0);
-    globalIdx.y = get_global_id(1);
-
-    if (!amount) {
-        if (globalIdx.x < width && globalIdx.y < height)
-            dst[globalIdx.x + globalIdx.y*dst_stride] = src[globalIdx.x + globalIdx.y*src_stride];
-        return;
-    }
-
-    local unsigned int l[32][32];
-    local unsigned int lcx[LU_RADIUS_X];
-    local unsigned int lcy[LU_RADIUS_Y];
-    int indexIx, indexIy, i, j;
-
-    //load up tile: actual workspace + halo of 8 points in x and y \n
-    for(i = 0; i <= 1; i++) {
-        indexIy = -8 + (blockIdx.y + i) * 16 + threadIdx.y;
-        indexIy = indexIy < 0 ? 0 : indexIy;
-        indexIy = indexIy >= height ? height - 1: indexIy;
-        for(j = 0; j <= 1; j++) {
-            indexIx = -8 + (blockIdx.x + j) * 16 + threadIdx.x;
-            indexIx = indexIx < 0 ? 0 : indexIx;
-            indexIx = indexIx >= width ? width - 1: indexIx;
-            l[i*16 + threadIdx.y][j*16 + threadIdx.x] = src[indexIy*src_stride + indexIx];
-        }
-    }
-
-    int indexL = threadIdx.y*16 + threadIdx.x;
-    if (indexL < LU_RADIUS_X)
-        lcx[indexL] = mask_x[indexL];
-    if (indexL < LU_RADIUS_Y)
-        lcy[indexL] = mask_y[indexL];
-    barrier(CLK_LOCAL_MEM_FENCE);
-
-    //needed for unsharp mask application in the end \n
-    int orig_value = (int)l[threadIdx.y + 8][threadIdx.x + 8];
-
-    int idx, idy, maskIndex;
-    int temp[2] = {0};
-    int steps_x = (LU_RADIUS_X-1)/2;
-    int steps_y = (LU_RADIUS_Y-1)/2;
-
-    // compute the actual workspace + left&right halos \n
-      \n#pragma unroll\n
-    for (j = 0; j <=1; j++) {
-      //extra work to cover left and right halos \n
-      idx = 16*j + threadIdx.x;
-      \n#pragma unroll\n
-        for (i = -steps_y; i <= steps_y; i++) {
-          idy = 8 + i + threadIdx.y;
-          maskIndex = (i + steps_y);
-          temp[j] += (int)l[idy][idx] * lcy[maskIndex];
-        }
-    }
-    barrier(CLK_LOCAL_MEM_FENCE);
-    //save results from the vertical filter in local memory \n
-    idy = 8 + threadIdx.y;
-      \n#pragma unroll\n
-    for (j = 0; j <=1; j++) {
-      idx = 16*j + threadIdx.x;
-      l[idy][idx] = temp[j];
-    }
-    barrier(CLK_LOCAL_MEM_FENCE);
-
-    //compute results with the horizontal filter \n
-    int sum = 0;
-    idy = 8 + threadIdx.y;
-    \n#pragma unroll\n
-      for (j = -steps_x; j <= steps_x; j++) {
-        idx = 8 + j + threadIdx.x;
-        maskIndex = j + steps_x;
-        sum += (int)l[idy][idx] * lcx[maskIndex];
-      }
-
-    int res = orig_value + (((orig_value - (int)((sum + halfscale) >> scalebits)) * amount) >> 16);
-
-    if (globalIdx.x < width && globalIdx.y < height)
-        dst[globalIdx.x + globalIdx.y*dst_stride] = clip_uint8(res);
-}
-
-kernel void unsharp_chroma(
-                    global unsigned char *src_y,
-                    global unsigned char *dst_y,
-                    global int *mask_x,
-                    global int *mask_y,
-                    int amount,
-                    int scalebits,
-                    int halfscale,
-                    int src_stride_lu,
-                    int src_stride_ch,
-                    int dst_stride_lu,
-                    int dst_stride_ch,
-                    int width,
-                    int height,
-                    int cw,
-                    int ch)
-{
-    global unsigned char *dst_u = dst_y + height * dst_stride_lu;
-    global unsigned char *dst_v = dst_u + ch * dst_stride_ch;
-    global unsigned char *src_u = src_y + height * src_stride_lu;
-    global unsigned char *src_v = src_u + ch * src_stride_ch;
-    int2 threadIdx, blockIdx, globalIdx;
-    threadIdx.x = get_local_id(0);
-    threadIdx.y = get_local_id(1);
-    blockIdx.x = get_group_id(0);
-    blockIdx.y = get_group_id(1);
-    globalIdx.x = get_global_id(0);
-    globalIdx.y = get_global_id(1);
-    int padch = get_global_size(1)/2;
-    global unsigned char *src = globalIdx.y>=padch ? src_v : src_u;
-    global unsigned char *dst = globalIdx.y>=padch ? dst_v : dst_u;
-
-    blockIdx.y = globalIdx.y>=padch ? blockIdx.y - get_num_groups(1)/2 : blockIdx.y;
-    globalIdx.y = globalIdx.y>=padch ? globalIdx.y - padch : globalIdx.y;
-
-    if (!amount) {
-        if (globalIdx.x < cw && globalIdx.y < ch)
-            dst[globalIdx.x + globalIdx.y*dst_stride_ch] = src[globalIdx.x + globalIdx.y*src_stride_ch];
-        return;
-    }
-
-    local unsigned int l[32][32];
-    local unsigned int lcx[CH_RADIUS_X];
-    local unsigned int lcy[CH_RADIUS_Y];
-    int indexIx, indexIy, i, j;
-    for(i = 0; i <= 1; i++) {
-        indexIy = -8 + (blockIdx.y + i) * 16 + threadIdx.y;
-        indexIy = indexIy < 0 ? 0 : indexIy;
-        indexIy = indexIy >= ch ? ch - 1: indexIy;
-        for(j = 0; j <= 1; j++) {
-            indexIx = -8 + (blockIdx.x + j) * 16 + threadIdx.x;
-            indexIx = indexIx < 0 ? 0 : indexIx;
-            indexIx = indexIx >= cw ? cw - 1: indexIx;
-            l[i*16 + threadIdx.y][j*16 + threadIdx.x] = src[indexIy * src_stride_ch + indexIx];
-        }
-    }
-
-    int indexL = threadIdx.y*16 + threadIdx.x;
-    if (indexL < CH_RADIUS_X)
-        lcx[indexL] = mask_x[indexL];
-    if (indexL < CH_RADIUS_Y)
-        lcy[indexL] = mask_y[indexL];
-    barrier(CLK_LOCAL_MEM_FENCE);
-
-    int orig_value = (int)l[threadIdx.y + 8][threadIdx.x + 8];
-
-    int idx, idy, maskIndex;
-    int steps_x = CH_RADIUS_X/2;
-    int steps_y = CH_RADIUS_Y/2;
-    int temp[2] = {0,0};
-
-    \n#pragma unroll\n
-      for (j = 0; j <= 1; j++) {
-        idx = 16*j + threadIdx.x;
-        \n#pragma unroll\n
-          for (i = -steps_y; i <= steps_y; i++) {
-            idy = 8 + i + threadIdx.y;
-            maskIndex = i + steps_y;
-            temp[j] += (int)l[idy][idx] * lcy[maskIndex];
-          }
-      }
-
-    barrier(CLK_LOCAL_MEM_FENCE);
-    idy = 8 + threadIdx.y;
-    \n#pragma unroll\n
-    for (j = 0; j <= 1; j++) {
-      idx = 16*j + threadIdx.x;
-      l[idy][idx] = temp[j];
-    }
-    barrier(CLK_LOCAL_MEM_FENCE);
-
-    //compute results with the horizontal filter \n
-    int sum = 0;
-    idy = 8 + threadIdx.y;
-    \n#pragma unroll\n
-      for (j = -steps_x; j <= steps_x; j++) {
-        idx = 8 + j + threadIdx.x;
-        maskIndex = j + steps_x;
-        sum += (int)l[idy][idx] * lcx[maskIndex];
-      }
-
-    int res = orig_value + (((orig_value - (int)((sum + halfscale) >> scalebits)) * amount) >> 16);
-
-    if (globalIdx.x < cw && globalIdx.y < ch)
-        dst[globalIdx.x + globalIdx.y*dst_stride_ch] = clip_uint8(res);
-}
-
-kernel void unsharp_default(global  unsigned char *src,
-                    global  unsigned char *dst,
-                    const global  unsigned int *mask_lu,
-                    const global  unsigned int *mask_ch,
-                    int amount_lu,
-                    int amount_ch,
-                    int step_x_lu,
-                    int step_y_lu,
-                    int step_x_ch,
-                    int step_y_ch,
-                    int scalebits_lu,
-                    int scalebits_ch,
-                    int halfscale_lu,
-                    int halfscale_ch,
-                    int src_stride_lu,
-                    int src_stride_ch,
-                    int dst_stride_lu,
-                    int dst_stride_ch,
-                    int height,
-                    int width,
-                    int ch,
-                    int cw)
-{
-    global unsigned char *dst_y = dst;
-    global unsigned char *dst_u = dst_y + height * dst_stride_lu;
-    global unsigned char *dst_v = dst_u + ch * dst_stride_ch;
-
-    global unsigned char *src_y = src;
-    global unsigned char *src_u = src_y + height * src_stride_lu;
-    global unsigned char *src_v = src_u + ch * src_stride_ch;
-
-    global unsigned char *temp_dst;
-    global unsigned char *temp_src;
-    const global unsigned int *temp_mask;
-    int global_id = get_global_id(0);
-    int i, j, x, y, temp_src_stride, temp_dst_stride, temp_height, temp_width, temp_steps_x, temp_steps_y,
-        temp_amount, temp_scalebits, temp_halfscale, sum, idx_x, idx_y, temp, res;
-    if (global_id < width * height) {
-        y = global_id / width;
-        x = global_id % width;
-        temp_dst = dst_y;
-        temp_src = src_y;
-        temp_src_stride = src_stride_lu;
-        temp_dst_stride = dst_stride_lu;
-        temp_height = height;
-        temp_width = width;
-        temp_steps_x = step_x_lu;
-        temp_steps_y = step_y_lu;
-        temp_mask = mask_lu;
-        temp_amount = amount_lu;
-        temp_scalebits = scalebits_lu;
-        temp_halfscale = halfscale_lu;
-    } else if ((global_id >= width * height) && (global_id < width * height + ch * cw)) {
-        y = (global_id - width * height) / cw;
-        x = (global_id - width * height) % cw;
-        temp_dst = dst_u;
-        temp_src = src_u;
-        temp_src_stride = src_stride_ch;
-        temp_dst_stride = dst_stride_ch;
-        temp_height = ch;
-        temp_width = cw;
-        temp_steps_x = step_x_ch;
-        temp_steps_y = step_y_ch;
-        temp_mask = mask_ch;
-        temp_amount = amount_ch;
-        temp_scalebits = scalebits_ch;
-        temp_halfscale = halfscale_ch;
-    } else {
-        y = (global_id - width * height - ch * cw) / cw;
-        x = (global_id - width * height - ch * cw) % cw;
-        temp_dst = dst_v;
-        temp_src = src_v;
-        temp_src_stride = src_stride_ch;
-        temp_dst_stride = dst_stride_ch;
-        temp_height = ch;
-        temp_width = cw;
-        temp_steps_x = step_x_ch;
-        temp_steps_y = step_y_ch;
-        temp_mask = mask_ch;
-        temp_amount = amount_ch;
-        temp_scalebits = scalebits_ch;
-        temp_halfscale = halfscale_ch;
-    }
-    if (temp_amount) {
-        sum = 0;
-        for (j = 0; j <= 2 * temp_steps_y; j++) {
-            idx_y = (y - temp_steps_y + j) <= 0 ? 0 : (y - temp_steps_y + j) >= temp_height ? temp_height-1 : y - temp_steps_y + j;
-            for (i = 0; i <= 2 * temp_steps_x; i++) {
-                idx_x = (x - temp_steps_x + i) <= 0 ? 0 : (x - temp_steps_x + i) >= temp_width ? temp_width-1 : x - temp_steps_x + i;
-                sum += temp_mask[i + j * (2 * temp_steps_x + 1)] * temp_src[idx_x + idx_y * temp_src_stride];
-            }
-        }
-        temp = (int)temp_src[x + y * temp_src_stride];
-        res = temp + (((temp - (int)((sum + temp_halfscale) >> temp_scalebits)) * temp_amount) >> 16);
-        temp_dst[x + y * temp_dst_stride] = clip_uint8(res);
-    } else {
-        temp_dst[x + y * temp_dst_stride] = temp_src[x + y * temp_src_stride];
-    }
-}
-);
-
-#endif /* AVFILTER_UNSHARP_OPENCL_KERNEL_H */
diff --git a/libavfilter/vf_deshake.c b/libavfilter/vf_deshake.c
index 64b48c6d02..fb4eb355b8 100644
--- a/libavfilter/vf_deshake.c
+++ b/libavfilter/vf_deshake.c
@@ -60,7 +60,6 @@
 #include "libavutil/qsort.h"
 
 #include "deshake.h"
-#include "deshake_opencl.h"
 
 #define OFFSET(x) offsetof(DeshakeContext, x)
 #define FLAGS AV_OPT_FLAG_VIDEO_PARAM|AV_OPT_FLAG_FILTERING_PARAM
@@ -83,7 +82,7 @@ static const AVOption deshake_options[] = {
         { "exhaustive", "exhaustive search",      0, AV_OPT_TYPE_CONST, {.i64=EXHAUSTIVE},       INT_MIN, INT_MAX, FLAGS, "smode" },
         { "less",       "less exhaustive search", 0, AV_OPT_TYPE_CONST, {.i64=SMART_EXHAUSTIVE}, INT_MIN, INT_MAX, FLAGS, "smode" },
     { "filename", "set motion search detailed log file name", OFFSET(filename), AV_OPT_TYPE_STRING, {.str=NULL}, .flags = FLAGS },
-    { "opencl", "use OpenCL filtering capabilities", OFFSET(opencl), AV_OPT_TYPE_BOOL, {.i64=0}, 0, 1, .flags = FLAGS },
+    { "opencl", "ignored",                              OFFSET(opencl), AV_OPT_TYPE_BOOL, {.i64=0}, 0, 1, .flags = FLAGS },
     { NULL }
 };
 
@@ -341,7 +340,6 @@ static int deshake_transform_c(AVFilterContext *ctx,
 
 static av_cold int init(AVFilterContext *ctx)
 {
-    int ret;
     DeshakeContext *deshake = ctx->priv;
 
     deshake->sad = av_pixelutils_get_sad_fn(4, 4, 1, deshake); // 16x16, 2nd source unaligned
@@ -369,17 +367,7 @@ static av_cold int init(AVFilterContext *ctx)
         deshake->cx &= ~15;
     }
     deshake->transform = deshake_transform_c;
-    if (!CONFIG_OPENCL && deshake->opencl) {
-        av_log(ctx, AV_LOG_ERROR, "OpenCL support was not enabled in this build, cannot be selected\n");
-        return AVERROR(EINVAL);
-    }
 
-    if (CONFIG_OPENCL && deshake->opencl) {
-        deshake->transform = ff_opencl_transform;
-        ret = ff_opencl_deshake_init(ctx);
-        if (ret < 0)
-            return ret;
-    }
     av_log(ctx, AV_LOG_VERBOSE, "cx: %d, cy: %d, cw: %d, ch: %d, rx: %d, ry: %d, edge: %d blocksize: %d contrast: %d search: %d\n",
            deshake->cx, deshake->cy, deshake->cw, deshake->ch,
            deshake->rx, deshake->ry, deshake->edge, deshake->blocksize * 2, deshake->contrast, deshake->search);
@@ -416,9 +404,6 @@ static int config_props(AVFilterLink *link)
 static av_cold void uninit(AVFilterContext *ctx)
 {
     DeshakeContext *deshake = ctx->priv;
-    if (CONFIG_OPENCL && deshake->opencl) {
-        ff_opencl_deshake_uninit(ctx);
-    }
     av_frame_free(&deshake->ref);
     av_freep(&deshake->angles);
     deshake->angles_size = 0;
@@ -447,12 +432,6 @@ static int filter_frame(AVFilterLink *link, AVFrame *in)
     }
     av_frame_copy_props(out, in);
 
-    if (CONFIG_OPENCL && deshake->opencl) {
-        ret = ff_opencl_deshake_process_inout_buf(link->dst,in, out);
-        if (ret < 0)
-            goto fail;
-    }
-
     if (deshake->cx < 0 || deshake->cy < 0 || deshake->cw < 0 || deshake->ch < 0) {
         // Find the most likely global motion for the current frame
         find_motion(deshake, (deshake->ref == NULL) ? in->data[0] : deshake->ref->data[0], in->data[0], link->w, link->h, in->linesize[0], &t);
diff --git a/libavfilter/vf_unsharp.c b/libavfilter/vf_unsharp.c
index 438ff6d8fb..41ccc56942 100644
--- a/libavfilter/vf_unsharp.c
+++ b/libavfilter/vf_unsharp.c
@@ -46,7 +46,6 @@
 #include "libavutil/opt.h"
 #include "libavutil/pixdesc.h"
 #include "unsharp.h"
-#include "unsharp_opencl.h"
 
 static void apply_unsharp(      uint8_t *dst, int dst_stride,
                           const uint8_t *src, int src_stride,
@@ -134,10 +133,8 @@ static void set_filter_param(UnsharpFilterParam *fp, int msize_x, int msize_y, f
 
 static av_cold int init(AVFilterContext *ctx)
 {
-    int ret = 0;
     UnsharpContext *s = ctx->priv;
 
-
     set_filter_param(&s->luma,   s->lmsize_x, s->lmsize_y, s->lamount);
     set_filter_param(&s->chroma, s->cmsize_x, s->cmsize_y, s->camount);
 
@@ -146,16 +143,6 @@ static av_cold int init(AVFilterContext *ctx)
         return AVERROR(EINVAL);
     }
     s->apply_unsharp = apply_unsharp_c;
-    if (!CONFIG_OPENCL && s->opencl) {
-        av_log(ctx, AV_LOG_ERROR, "OpenCL support was not enabled in this build, cannot be selected\n");
-        return AVERROR(EINVAL);
-    }
-    if (CONFIG_OPENCL && s->opencl) {
-        s->apply_unsharp = ff_opencl_apply_unsharp;
-        ret = ff_opencl_unsharp_init(ctx);
-        if (ret < 0)
-            return ret;
-    }
     return 0;
 }
 
@@ -227,10 +214,6 @@ static av_cold void uninit(AVFilterContext *ctx)
 {
     UnsharpContext *s = ctx->priv;
 
-    if (CONFIG_OPENCL && s->opencl) {
-        ff_opencl_unsharp_uninit(ctx);
-    }
-
     free_filter_param(&s->luma);
     free_filter_param(&s->chroma);
 }
@@ -248,14 +231,9 @@ static int filter_frame(AVFilterLink *link, AVFrame *in)
         return AVERROR(ENOMEM);
     }
     av_frame_copy_props(out, in);
-    if (CONFIG_OPENCL && s->opencl) {
-        ret = ff_opencl_unsharp_process_inout_buf(link->dst, in, out);
-        if (ret < 0)
-            goto end;
-    }
 
     ret = s->apply_unsharp(link->dst, in, out);
-end:
+
     av_frame_free(&in);
 
     if (ret < 0) {
@@ -282,7 +260,7 @@ static const AVOption unsharp_options[] = {
     { "cy",             "set chroma matrix vertical size",   OFFSET(cmsize_y), AV_OPT_TYPE_INT,   { .i64 = 5 }, MIN_SIZE, MAX_SIZE, FLAGS },
     { "chroma_amount",  "set chroma effect strength",        OFFSET(camount),  AV_OPT_TYPE_FLOAT, { .dbl = 0 },       -2,        5, FLAGS },
     { "ca",             "set chroma effect strength",        OFFSET(camount),  AV_OPT_TYPE_FLOAT, { .dbl = 0 },       -2,        5, FLAGS },
-    { "opencl",         "use OpenCL filtering capabilities", OFFSET(opencl),   AV_OPT_TYPE_BOOL,  { .i64 = 0 },        0,        1, FLAGS },
+    { "opencl",         "ignored",                           OFFSET(opencl),   AV_OPT_TYPE_BOOL,  { .i64 = 0 },        0,        1, FLAGS },
     { NULL }
 };