avfilter/opencl: compile kernels separately

Reviewed-by: Wei Gao <highgod0401@gmail.com>
Signed-off-by: Michael Niedermayer <michaelni@gmx.at>

5 files changed, 51 insertions, 24 deletions

diff --git a/libavfilter/deshake.h b/libavfilter/deshake.h
index c24090ee36..5792973957 100644
--- a/libavfilter/deshake.h
+++ b/libavfilter/deshake.h
@@ -55,6 +55,9 @@ typedef struct {
 #if CONFIG_OPENCL
 
 typedef struct {
+    cl_command_queue command_queue;
+    cl_program program;
+    cl_kernel kernel;
     size_t matrix_size;
     float matrix_y[9];
     float matrix_uv[9];
@@ -67,7 +70,6 @@ typedef struct {
     size_t cl_inbuf_size;
     cl_mem cl_outbuf;
     size_t cl_outbuf_size;
-    AVOpenCLKernelEnv kernel_env;
 } DeshakeOpenclContext;
 
 #endif
diff --git a/libavfilter/deshake_opencl.c b/libavfilter/deshake_opencl.c
index eea873ea4a..e4e4df19e8 100644
--- a/libavfilter/deshake_opencl.c
+++ b/libavfilter/deshake_opencl.c
@@ -45,7 +45,7 @@ int ff_opencl_transform(AVFilterContext *ctx,
     FFOpenclParam opencl_param = {0};
 
     opencl_param.ctx = ctx;
-    opencl_param.kernel = deshake->opencl_ctx.kernel_env.kernel;
+    opencl_param.kernel = deshake->opencl_ctx.kernel;
     ret = av_opencl_buffer_write(deshake->opencl_ctx.cl_matrix_y, (uint8_t *)matrix_y, deshake->opencl_ctx.matrix_size * sizeof(cl_float));
     if (ret < 0)
         return ret;
@@ -75,14 +75,14 @@ int ff_opencl_transform(AVFilterContext *ctx,
                                   NULL);
     if (ret < 0)
         return ret;
-    status = clEnqueueNDRangeKernel(deshake->opencl_ctx.kernel_env.command_queue,
-                                    deshake->opencl_ctx.kernel_env.kernel, 1, NULL,
+    status = clEnqueueNDRangeKernel(deshake->opencl_ctx.command_queue,
+                                    deshake->opencl_ctx.kernel, 1, NULL,
                                     &global_work_size, 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;
     }
-    clFinish(deshake->opencl_ctx.kernel_env.command_queue);
+    clFinish(deshake->opencl_ctx.command_queue);
     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);
@@ -108,11 +108,21 @@ int ff_opencl_deshake_init(AVFilterContext *ctx)
         deshake->opencl_ctx.matrix_size*sizeof(cl_float), CL_MEM_READ_ONLY, NULL);
     if (ret < 0)
         return ret;
-    if (!deshake->opencl_ctx.kernel_env.kernel) {
-        ret =  av_opencl_create_kernel(&deshake->opencl_ctx.kernel_env, "avfilter_transform");
-        if (ret < 0) {
-            av_log(ctx, AV_LOG_ERROR, "OpenCL failed to create kernel for name 'avfilter_transform'\n");
-            return ret;
+    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) {
+        deshake->opencl_ctx.kernel = clCreateKernel(deshake->opencl_ctx.program, "avfilter_transform", &ret);
+        if (ret != CL_SUCCESS) {
+            av_log(ctx, AV_LOG_ERROR, "OpenCL failed to create kernel 'avfilter_transform'\n");
+            return AVERROR(EINVAL);
         }
     }
     return ret;
@@ -125,11 +135,12 @@ void ff_opencl_deshake_uninit(AVFilterContext *ctx)
     av_opencl_buffer_release(&deshake->opencl_ctx.cl_outbuf);
     av_opencl_buffer_release(&deshake->opencl_ctx.cl_matrix_y);
     av_opencl_buffer_release(&deshake->opencl_ctx.cl_matrix_uv);
-    av_opencl_release_kernel(&deshake->opencl_ctx.kernel_env);
+    clReleaseKernel(deshake->opencl_ctx.kernel);
+    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;
diff --git a/libavfilter/unsharp.h b/libavfilter/unsharp.h
index c2259298f6..2738243cea 100644
--- a/libavfilter/unsharp.h
+++ b/libavfilter/unsharp.h
@@ -33,6 +33,9 @@
 #if CONFIG_OPENCL
 
 typedef struct {
+    cl_command_queue command_queue;
+    cl_program program;
+    cl_kernel kernel;
     cl_mem cl_luma_mask;
     cl_mem cl_chroma_mask;
     int in_plane_size[8];
@@ -42,7 +45,6 @@ typedef struct {
     size_t cl_inbuf_size;
     cl_mem cl_outbuf;
     size_t cl_outbuf_size;
-    AVOpenCLKernelEnv kernel_env;
 } UnsharpOpenclContext;
 
 #endif
diff --git a/libavfilter/unsharp_opencl.c b/libavfilter/unsharp_opencl.c
index b373b66116..77b59371a8 100644
--- a/libavfilter/unsharp_opencl.c
+++ b/libavfilter/unsharp_opencl.c
@@ -159,7 +159,7 @@ int ff_opencl_apply_unsharp(AVFilterContext *ctx, AVFrame *in, AVFrame *out)
     FFOpenclParam opencl_param = {0};
 
     opencl_param.ctx = ctx;
-    opencl_param.kernel = unsharp->opencl_ctx.kernel_env.kernel;
+    opencl_param.kernel = unsharp->opencl_ctx.kernel;
     ret = ff_opencl_set_parameter(&opencl_param,
                                   FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_inbuf),
                                   FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_outbuf),
@@ -186,14 +186,14 @@ int ff_opencl_apply_unsharp(AVFilterContext *ctx, AVFrame *in, AVFrame *out)
                                   NULL);
     if (ret < 0)
         return ret;
-    status = clEnqueueNDRangeKernel(unsharp->opencl_ctx.kernel_env.command_queue,
-                                    unsharp->opencl_ctx.kernel_env.kernel, 1, NULL,
+    status = clEnqueueNDRangeKernel(unsharp->opencl_ctx.command_queue,
+                                    unsharp->opencl_ctx.kernel, 1, NULL,
                                     &global_work_size, 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;
     }
-    clFinish(unsharp->opencl_ctx.kernel_env.command_queue);
+    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);
@@ -220,11 +220,21 @@ int ff_opencl_unsharp_init(AVFilterContext *ctx)
     if (ret < 0)
         return ret;
     unsharp->opencl_ctx.plane_num = PLANE_NUM;
-    if (!unsharp->opencl_ctx.kernel_env.kernel) {
-        ret = av_opencl_create_kernel(&unsharp->opencl_ctx.kernel_env, "unsharp");
-        if (ret < 0) {
-            av_log(ctx, AV_LOG_ERROR, "OpenCL failed to create kernel with name 'unsharp'\n");
-            return ret;
+    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);
+    }
+    unsharp->opencl_ctx.program = av_opencl_compile("unsharp", NULL);
+    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.kernel) {
+        unsharp->opencl_ctx.kernel = clCreateKernel(unsharp->opencl_ctx.program, "unsharp", &ret);
+        if (ret != CL_SUCCESS) {
+            av_log(ctx, AV_LOG_ERROR, "OpenCL failed to create kernel 'unsharp'\n");
+            return AVERROR(EINVAL);
         }
     }
     return ret;
@@ -237,7 +247,9 @@ void ff_opencl_unsharp_uninit(AVFilterContext *ctx)
     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_release_kernel(&unsharp->opencl_ctx.kernel_env);
+    clReleaseKernel(unsharp->opencl_ctx.kernel);
+    clReleaseProgram(unsharp->opencl_ctx.program);
+    unsharp->opencl_ctx.command_queue = NULL;
     av_opencl_uninit();
 }
 
diff --git a/libavfilter/version.h b/libavfilter/version.h
index a293d3bfd2..546ab5b338 100644
--- a/libavfilter/version.h
+++ b/libavfilter/version.h
@@ -31,7 +31,7 @@
 
 #define LIBAVFILTER_VERSION_MAJOR  3
 #define LIBAVFILTER_VERSION_MINOR  90
-#define LIBAVFILTER_VERSION_MICRO 101
+#define LIBAVFILTER_VERSION_MICRO 102
 
 #define LIBAVFILTER_VERSION_INT AV_VERSION_INT(LIBAVFILTER_VERSION_MAJOR, \
                                                LIBAVFILTER_VERSION_MINOR, \