[FFmpeg-devel] [PATCH] avfilter/vf_program_opencl: allow setting kernel per plane

Paul B Mahol onemda at gmail.com
Mon Feb 24 12:01:21 EET 2020


Fixes #7190

Signed-off-by: Paul B Mahol <onemda at gmail.com>
---
 doc/filters.texi                | 22 ++++++++++++
 libavfilter/vf_program_opencl.c | 64 ++++++++++++++++++++++-----------
 2 files changed, 65 insertions(+), 21 deletions(-)

diff --git a/doc/filters.texi b/doc/filters.texi
index 70fd7a4cc7..6b10f649b9 100644
--- a/doc/filters.texi
+++ b/doc/filters.texi
@@ -21302,6 +21302,17 @@ Number of inputs to the filter.  Defaults to 1.
 @item size, s
 Size of output frames.  Defaults to the same as the first input.
 
+ at item kernel2
+Kernel name in program for 2nd plane, if not set kernel from option
+ at var{kernel} is used.
+
+ at item kernel3
+Kernel name in program for 3nd plane, if not set kernel from option
+ at var{kernel} is used.
+
+ at item kernel4
+Kernel name in program for 4nd plane, if not set kernel from option
+ at var{kernel} is used.
 @end table
 
 The program source file must contain a kernel function with the given name,
@@ -22488,6 +22499,17 @@ Pixel format to use for the generated frames.  This must be set.
 @item rate, r
 Number of frames generated every second.  Default value is '25'.
 
+ at item kernel2
+Kernel name in program for 2nd plane, if not set kernel from option
+ at var{kernel} is used.
+
+ at item kernel3
+Kernel name in program for 3nd plane, if not set kernel from option
+ at var{kernel} is used.
+
+ at item kernel4
+Kernel name in program for 4nd plane, if not set kernel from option
+ at var{kernel} is used.
 @end table
 
 For details of how the program loading works, see the @ref{program_opencl}
diff --git a/libavfilter/vf_program_opencl.c b/libavfilter/vf_program_opencl.c
index ec25e931f5..f748b15037 100644
--- a/libavfilter/vf_program_opencl.c
+++ b/libavfilter/vf_program_opencl.c
@@ -33,14 +33,14 @@ typedef struct ProgramOpenCLContext {
 
     int                 loaded;
     cl_uint             index;
-    cl_kernel           kernel;
+    cl_kernel           kernel[4];
     cl_command_queue    command_queue;
 
     FFFrameSync         fs;
     AVFrame           **frames;
 
     const char         *source_file;
-    const char         *kernel_name;
+    const char         *kernel_name[4];
     int                 nb_inputs;
     int                 width, height;
     enum AVPixelFormat  source_format;
@@ -66,15 +66,17 @@ static int program_opencl_load(AVFilterContext *avctx)
         return AVERROR(EIO);
     }
 
-    ctx->kernel = clCreateKernel(ctx->ocf.program, ctx->kernel_name, &cle);
-    if (!ctx->kernel) {
-        if (cle == CL_INVALID_KERNEL_NAME) {
-            av_log(avctx, AV_LOG_ERROR, "Kernel function '%s' not found in "
-                   "program.\n", ctx->kernel_name);
-        } else {
-            av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle);
+    for (int i = 0; i < 4; i++) {
+        ctx->kernel[i] = clCreateKernel(ctx->ocf.program, ctx->kernel_name[i] ? ctx->kernel_name[i] : ctx->kernel_name[0], &cle);
+        if (!ctx->kernel[i]) {
+            if (cle == CL_INVALID_KERNEL_NAME) {
+                av_log(avctx, AV_LOG_ERROR, "Kernel function '%s' not found in "
+                       "program.\n", ctx->kernel_name[i] ? ctx->kernel_name[i] : ctx->kernel_name[0]);
+            } else {
+                av_log(avctx, AV_LOG_ERROR, "Failed to create kernel%d: %d.\n", i, cle);
+            }
+            return AVERROR(EIO);
         }
-        return AVERROR(EIO);
     }
 
     ctx->loaded = 1;
@@ -108,14 +110,14 @@ static int program_opencl_run(AVFilterContext *avctx)
         if (!dst)
             break;
 
-        cle = clSetKernelArg(ctx->kernel, 0, sizeof(cl_mem), &dst);
+        cle = clSetKernelArg(ctx->kernel[plane], 0, sizeof(cl_mem), &dst);
         if (cle != CL_SUCCESS) {
             av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
                    "destination image argument: %d.\n", cle);
             err = AVERROR_UNKNOWN;
             goto fail;
         }
-        cle = clSetKernelArg(ctx->kernel, 1, sizeof(cl_uint), &ctx->index);
+        cle = clSetKernelArg(ctx->kernel[plane], 1, sizeof(cl_uint), &ctx->index);
         if (cle != CL_SUCCESS) {
             av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
                    "index argument: %d.\n", cle);
@@ -129,7 +131,7 @@ static int program_opencl_run(AVFilterContext *avctx)
             src = (cl_mem)ctx->frames[input]->data[plane];
             av_assert0(src);
 
-            cle = clSetKernelArg(ctx->kernel, 2 + input, sizeof(cl_mem), &src);
+            cle = clSetKernelArg(ctx->kernel[plane], 2 + input, sizeof(cl_mem), &src);
             if (cle != CL_SUCCESS) {
                 av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
                        "source image argument %d: %d.\n", input, cle);
@@ -147,7 +149,7 @@ static int program_opencl_run(AVFilterContext *avctx)
                "(%"SIZE_SPECIFIER"x%"SIZE_SPECIFIER").\n",
                plane, global_work[0], global_work[1]);
 
-        cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
+        cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel[plane], 2, NULL,
                                      global_work, NULL, 0, NULL, NULL);
         CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue kernel: %d.\n", cle);
     }
@@ -312,11 +314,13 @@ static av_cold void program_opencl_uninit(AVFilterContext *avctx)
             av_freep(&avctx->input_pads[i].name);
     }
 
-    if (ctx->kernel) {
-        cle = clReleaseKernel(ctx->kernel);
-        if (cle != CL_SUCCESS)
-            av_log(avctx, AV_LOG_ERROR, "Failed to release "
-                   "kernel: %d.\n", cle);
+    for (i = 0; i < 4; i++) {
+        if (ctx->kernel[i]) {
+            cle = clReleaseKernel(ctx->kernel[i]);
+            if (cle != CL_SUCCESS)
+                av_log(avctx, AV_LOG_ERROR, "Failed to release "
+                       "kernel%d: %d.\n", i, cle);
+        }
     }
 
     if (ctx->command_queue) {
@@ -337,7 +341,7 @@ static av_cold void program_opencl_uninit(AVFilterContext *avctx)
 static const AVOption program_opencl_options[] = {
     { "source", "OpenCL program source file", OFFSET(source_file),
       AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS },
-    { "kernel", "Kernel name in program",     OFFSET(kernel_name),
+    { "kernel", "Kernel name in program",     OFFSET(kernel_name[0]),
       AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS },
 
     { "inputs", "Number of inputs", OFFSET(nb_inputs),
@@ -348,6 +352,15 @@ static const AVOption program_opencl_options[] = {
     { "s",      "Video size",       OFFSET(width),
       AV_OPT_TYPE_IMAGE_SIZE,       { .str = NULL }, 0, 0, FLAGS },
 
+    { "kernel2", "Kernel name in program for 2nd plane", OFFSET(kernel_name[1]),
+      AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS },
+
+    { "kernel3", "Kernel name in program for 3rd plane", OFFSET(kernel_name[2]),
+      AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS },
+
+    { "kernel4", "Kernel name in program for 4th plane", OFFSET(kernel_name[3]),
+      AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS },
+
     { NULL },
 };
 
@@ -384,7 +397,7 @@ AVFilter ff_vf_program_opencl = {
 static const AVOption openclsrc_options[] = {
     { "source", "OpenCL program source file", OFFSET(source_file),
       AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS },
-    { "kernel", "Kernel name in program",     OFFSET(kernel_name),
+    { "kernel", "Kernel name in program",     OFFSET(kernel_name[0]),
       AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS },
 
     { "size",   "Video size",       OFFSET(width),
@@ -400,6 +413,15 @@ static const AVOption openclsrc_options[] = {
     { "r",      "Video frame rate", OFFSET(source_rate),
       AV_OPT_TYPE_VIDEO_RATE,       { .str = "25" }, 0, INT_MAX, FLAGS },
 
+    { "kernel2", "Kernel name in program for 2nd plane", OFFSET(kernel_name[1]),
+      AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS },
+
+    { "kernel3", "Kernel name in program for 3rd plane", OFFSET(kernel_name[2]),
+      AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS },
+
+    { "kernel4", "Kernel name in program for 4th plane", OFFSET(kernel_name[3]),
+      AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS },
+
     { NULL },
 };
 
-- 
2.17.1



More information about the ffmpeg-devel mailing list