[FFmpeg-devel,1/2] lavfi/opencl: add macro for opencl error handling.

Submitted by Ruiling Song on July 2, 2018, 8:55 p.m.

Details

Message ID 1530564953-4747-1-git-send-email-ruiling.song@intel.com
State New
Headers show

Commit Message

Ruiling Song July 2, 2018, 8:55 p.m.
Signed-off-by: Ruiling Song <ruiling.song@intel.com>
---
 libavfilter/opencl.h            |  4 ++--
 libavfilter/vf_avgblur_opencl.c | 45 +++++++++--------------------------
 libavfilter/vf_overlay_opencl.c | 29 +++++------------------
 libavfilter/vf_program_opencl.c | 14 ++---------
 libavfilter/vf_tonemap_opencl.c | 33 +++++---------------------
 libavfilter/vf_unsharp_opencl.c | 52 +++++++++--------------------------------
 6 files changed, 38 insertions(+), 139 deletions(-)

Comments

Mark Thompson July 2, 2018, 10:40 p.m.
On 02/07/18 21:55, Ruiling Song wrote:
> Signed-off-by: Ruiling Song <ruiling.song@intel.com>
> ---
>  libavfilter/opencl.h            |  4 ++--
>  libavfilter/vf_avgblur_opencl.c | 45 +++++++++--------------------------
>  libavfilter/vf_overlay_opencl.c | 29 +++++------------------
>  libavfilter/vf_program_opencl.c | 14 ++---------
>  libavfilter/vf_tonemap_opencl.c | 33 +++++---------------------
>  libavfilter/vf_unsharp_opencl.c | 52 +++++++++--------------------------------
>  6 files changed, 38 insertions(+), 139 deletions(-)

This doesn't apply - it seems like it's on top of another patch?

Code looks good to me.

Thanks,

- Mark


> diff --git a/libavfilter/opencl.h b/libavfilter/opencl.h
> index fd76f72..0ed360b 100644
> --- a/libavfilter/opencl.h
> +++ b/libavfilter/opencl.h
> @@ -116,9 +116,9 @@ int ff_opencl_filter_work_size_from_image(AVFilterContext *avctx,
>   * A helper macro to handle OpenCL error. It will assign errcode to
>   * variable err, log error msg, and jump to fail label on error.
>   */
> -#define OCL_FAIL_ON_ERR(logctx, cle, errcode, ...) do {\
> +#define CL_FAIL_ON_ERROR(errcode, ...) do {\
>      if (cle != CL_SUCCESS) {\
> -        av_log(logctx, AV_LOG_ERROR, __VA_ARGS__);\
> +        av_log(avctx, AV_LOG_ERROR, __VA_ARGS__);\
>          err = errcode;\
>          goto fail;\
>      }\
> diff --git a/libavfilter/vf_avgblur_opencl.c b/libavfilter/vf_avgblur_opencl.c
> index d1d3eb1..bc6bcab 100644
> --- a/libavfilter/vf_avgblur_opencl.c
> +++ b/libavfilter/vf_avgblur_opencl.c
> @@ -64,26 +64,16 @@ static int avgblur_opencl_init(AVFilterContext *avctx)
>      ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
>                                                ctx->ocf.hwctx->device_id,
>                                                0, &cle);
> -    if (!ctx->command_queue) {
> -        av_log(avctx, AV_LOG_ERROR, "Failed to create OpenCL "
> -               "command queue: %d.\n", cle);
> -        err = AVERROR(EIO);
> -        goto fail;
> -    }
> +    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL "
> +                     "command queue %d.\n", cle);
>  
>      ctx->kernel_horiz = clCreateKernel(ctx->ocf.program,"avgblur_horiz", &cle);
> -    if (!ctx->kernel_horiz) {
> -        av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle);
> -        err = AVERROR(EIO);
> -        goto fail;
> -    }
> +    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create horizontal "
> +                     "kernel %d.\n", cle);
>  
>      ctx->kernel_vert = clCreateKernel(ctx->ocf.program,"avgblur_vert", &cle);
> -    if (!ctx->kernel_vert) {
> -        av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle);
> -        err = AVERROR(EIO);
> -        goto fail;
> -    }
> +    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create vertical "
> +                     "kernel %d.\n", cle);
>  
>      ctx->initialised = 1;
>      return 0;
> @@ -236,12 +226,8 @@ static int avgblur_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
>              cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel_horiz, 2, NULL,
>                                           global_work, NULL,
>                                           0, NULL, NULL);
> -            if (cle != CL_SUCCESS) {
> -                av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n",
> -                       cle);
> -                err = AVERROR(EIO);
> -                goto fail;
> -            }
> +            CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue horizontal "
> +                             "kernel: %d.\n", cle);
>              cle = clFinish(ctx->command_queue);
>  
>              err = ff_opencl_filter_work_size_from_image(avctx, global_work,
> @@ -259,22 +245,13 @@ static int avgblur_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
>              cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel_vert, 2, NULL,
>                                           global_work, NULL,
>                                           0, NULL, NULL);
> -            if (cle != CL_SUCCESS) {
> -                av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n",
> -                       cle);
> -                err = AVERROR(EIO);
> -                goto fail;
> -            }
> +            CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue vertical "
> +                             "kernel: %d.\n", cle);
>          }
>      }
>  
>      cle = clFinish(ctx->command_queue);
> -    if (cle != CL_SUCCESS) {
> -        av_log(avctx, AV_LOG_ERROR, "Failed to finish command queue: %d.\n",
> -               cle);
> -        err = AVERROR(EIO);
> -        goto fail;
> -    }
> +    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
>  
>      err = av_frame_copy_props(output, input);
>      if (err < 0)
> diff --git a/libavfilter/vf_overlay_opencl.c b/libavfilter/vf_overlay_opencl.c
> index 556ce35..e9c8532 100644
> --- a/libavfilter/vf_overlay_opencl.c
> +++ b/libavfilter/vf_overlay_opencl.c
> @@ -100,19 +100,11 @@ static int overlay_opencl_load(AVFilterContext *avctx,
>      ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
>                                                ctx->ocf.hwctx->device_id,
>                                                0, &cle);
> -    if (!ctx->command_queue) {
> -        av_log(avctx, AV_LOG_ERROR, "Failed to create OpenCL "
> -               "command queue: %d.\n", cle);
> -        err = AVERROR(EIO);
> -        goto fail;
> -    }
> +    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL "
> +                     "command queue %d.\n", cle);
>  
>      ctx->kernel = clCreateKernel(ctx->ocf.program, kernel, &cle);
> -    if (!ctx->kernel) {
> -        av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle);
> -        err = AVERROR(EIO);
> -        goto fail;
> -    }
> +    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle);
>  
>      ctx->initialised = 1;
>      return 0;
> @@ -209,21 +201,12 @@ static int overlay_opencl_blend(FFFrameSync *fs)
>  
>          cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
>                                       global_work, NULL, 0, NULL, NULL);
> -        if (cle != CL_SUCCESS) {
> -            av_log(avctx, AV_LOG_ERROR, "Failed to enqueue "
> -                   "overlay kernel for plane %d: %d.\n", cle, plane);
> -            err = AVERROR(EIO);
> -            goto fail;
> -        }
> +        CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue overlay kernel "
> +                         "for plane %d: %d.\n", plane, cle);
>      }
>  
>      cle = clFinish(ctx->command_queue);
> -    if (cle != CL_SUCCESS) {
> -        av_log(avctx, AV_LOG_ERROR, "Failed to finish "
> -               "command queue: %d.\n", cle);
> -        err = AVERROR(EIO);
> -        goto fail;
> -    }
> +    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
>  
>      err = av_frame_copy_props(output, input_main);
>  
> diff --git a/libavfilter/vf_program_opencl.c b/libavfilter/vf_program_opencl.c
> index a002792..dfb2565 100644
> --- a/libavfilter/vf_program_opencl.c
> +++ b/libavfilter/vf_program_opencl.c
> @@ -148,21 +148,11 @@ static int program_opencl_run(AVFilterContext *avctx)
>  
>          cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
>                                       global_work, NULL, 0, NULL, NULL);
> -        if (cle != CL_SUCCESS) {
> -            av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n",
> -                   cle);
> -            err = AVERROR(EIO);
> -            goto fail;
> -        }
> +        CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue kernel: %d.\n", cle);
>      }
>  
>      cle = clFinish(ctx->command_queue);
> -    if (cle != CL_SUCCESS) {
> -        av_log(avctx, AV_LOG_ERROR, "Failed to finish command queue: %d.\n",
> -               cle);
> -        err = AVERROR(EIO);
> -        goto fail;
> -    }
> +    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
>  
>      if (ctx->nb_inputs > 0) {
>          err = av_frame_copy_props(output, ctx->frames[0]);
> diff --git a/libavfilter/vf_tonemap_opencl.c b/libavfilter/vf_tonemap_opencl.c
> index 36c7fbe..241f95e 100644
> --- a/libavfilter/vf_tonemap_opencl.c
> +++ b/libavfilter/vf_tonemap_opencl.c
> @@ -262,29 +262,17 @@ static int tonemap_opencl_init(AVFilterContext *avctx)
>      ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
>                                                ctx->ocf.hwctx->device_id,
>                                                0, &cle);
> -    if (!ctx->command_queue) {
> -        av_log(avctx, AV_LOG_ERROR, "Failed to create OpenCL "
> -               "command queue: %d.\n", cle);
> -        err = AVERROR(EIO);
> -        goto fail;
> -    }
> +    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL "
> +                     "command queue %d.\n", cle);
>  
>      ctx->kernel = clCreateKernel(ctx->ocf.program, "tonemap", &cle);
> -    if (!ctx->kernel) {
> -        av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle);
> -        err = AVERROR(EIO);
> -        goto fail;
> -    }
> +    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle);
>  
>      ctx->util_mem =
>          clCreateBuffer(ctx->ocf.hwctx->context, 0,
>                         (2 * DETECTION_FRAMES + 7) * sizeof(unsigned),
>                         NULL, &cle);
> -    if (cle != CL_SUCCESS) {
> -        av_log(avctx, AV_LOG_ERROR, "Failed to create util buffer: %d.\n", cle);
> -        err = AVERROR(EIO);
> -        goto fail;
> -    }
> +    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create util buffer: %d.\n", cle);
>  
>      ctx->initialised = 1;
>      return 0;
> @@ -349,11 +337,7 @@ static int launch_kernel(AVFilterContext *avctx, cl_kernel kernel,
>      cle = clEnqueueNDRangeKernel(ctx->command_queue, kernel, 2, NULL,
>                                   global_work, local_work,
>                                   0, NULL, NULL);
> -    if (cle != CL_SUCCESS) {
> -        av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n",
> -               cle);
> -        return AVERROR(EIO);
> -    }
> +    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue kernel: %d.\n", cle);
>      return 0;
>  fail:
>      return err;
> @@ -482,12 +466,7 @@ static int tonemap_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
>      }
>  
>      cle = clFinish(ctx->command_queue);
> -    if (cle != CL_SUCCESS) {
> -        av_log(avctx, AV_LOG_ERROR, "Failed to finish command queue: %d.\n",
> -               cle);
> -        err = AVERROR(EIO);
> -        goto fail;
> -    }
> +    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
>  
>      av_frame_free(&input);
>  
> diff --git a/libavfilter/vf_unsharp_opencl.c b/libavfilter/vf_unsharp_opencl.c
> index 5b1eb59..d76d1b1 100644
> --- a/libavfilter/vf_unsharp_opencl.c
> +++ b/libavfilter/vf_unsharp_opencl.c
> @@ -76,12 +76,8 @@ static int unsharp_opencl_init(AVFilterContext *avctx)
>      ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
>                                                ctx->ocf.hwctx->device_id,
>                                                0, &cle);
> -    if (!ctx->command_queue) {
> -        av_log(avctx, AV_LOG_ERROR, "Failed to create OpenCL "
> -               "command queue: %d.\n", cle);
> -        err = AVERROR(EIO);
> -        goto fail;
> -    }
> +    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL "
> +                     "command queue %d.\n", cle);
>  
>      // Use global kernel if mask size will be too big for the local store..
>      ctx->global = (ctx->luma_size_x   > 17.0f ||
> @@ -92,11 +88,7 @@ static int unsharp_opencl_init(AVFilterContext *avctx)
>      ctx->kernel = clCreateKernel(ctx->ocf.program,
>                                   ctx->global ? "unsharp_global"
>                                               : "unsharp_local", &cle);
> -    if (!ctx->kernel) {
> -        av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle);
> -        err = AVERROR(EIO);
> -        goto fail;
> -    }
> +    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle);
>  
>      ctx->initialised = 1;
>      return 0;
> @@ -176,12 +168,8 @@ static int unsharp_opencl_make_filter_params(AVFilterContext *avctx)
>                                      CL_MEM_COPY_HOST_PTR |
>                                      CL_MEM_HOST_NO_ACCESS,
>                                      matrix_bytes, matrix, &cle);
> -            if (!buffer) {
> -                av_log(avctx, AV_LOG_ERROR, "Failed to create matrix buffer: "
> -                       "%d.\n", cle);
> -                err = AVERROR(EIO);
> -                goto fail;
> -            }
> +            CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create matrix buffer: "
> +                             "%d.\n", cle);
>              ctx->plane[p].matrix = buffer;
>          } else {
>              buffer = clCreateBuffer(ctx->ocf.hwctx->context,
> @@ -190,12 +178,8 @@ static int unsharp_opencl_make_filter_params(AVFilterContext *avctx)
>                                      CL_MEM_HOST_NO_ACCESS,
>                                      sizeof(ctx->plane[p].blur_x),
>                                      ctx->plane[p].blur_x, &cle);
> -            if (!buffer) {
> -                av_log(avctx, AV_LOG_ERROR, "Failed to create x-coef buffer: "
> -                       "%d.\n", cle);
> -                err = AVERROR(EIO);
> -                goto fail;
> -            }
> +            CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create x-coef buffer: "
> +                             "%d.\n", cle);
>              ctx->plane[p].coef_x = buffer;
>  
>              buffer = clCreateBuffer(ctx->ocf.hwctx->context,
> @@ -204,12 +188,8 @@ static int unsharp_opencl_make_filter_params(AVFilterContext *avctx)
>                                      CL_MEM_HOST_NO_ACCESS,
>                                      sizeof(ctx->plane[p].blur_y),
>                                      ctx->plane[p].blur_y, &cle);
> -            if (!buffer) {
> -                av_log(avctx, AV_LOG_ERROR, "Failed to create y-coef buffer: "
> -                       "%d.\n", cle);
> -                err = AVERROR(EIO);
> -                goto fail;
> -            }
> +            CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create y-coef buffer: "
> +                             "%d.\n", cle);
>              ctx->plane[p].coef_y = buffer;
>          }
>  
> @@ -296,21 +276,11 @@ static int unsharp_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
>          cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
>                                       global_work, ctx->global ? NULL : local_work,
>                                       0, NULL, NULL);
> -        if (cle != CL_SUCCESS) {
> -            av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n",
> -                   cle);
> -            err = AVERROR(EIO);
> -            goto fail;
> -        }
> +        CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue kernel: %d.\n", cle);
>      }
>  
>      cle = clFinish(ctx->command_queue);
> -    if (cle != CL_SUCCESS) {
> -        av_log(avctx, AV_LOG_ERROR, "Failed to finish command queue: %d.\n",
> -               cle);
> -        err = AVERROR(EIO);
> -        goto fail;
> -    }
> +    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
>  
>      err = av_frame_copy_props(output, input);
>      if (err < 0)
>
Ruiling Song July 3, 2018, 6:19 p.m.
> -----Original Message-----

> From: ffmpeg-devel [mailto:ffmpeg-devel-bounces@ffmpeg.org] On Behalf Of

> Mark Thompson

> Sent: Monday, July 2, 2018 3:40 PM

> To: ffmpeg-devel@ffmpeg.org

> Subject: Re: [FFmpeg-devel] [PATCH 1/2] lavfi/opencl: add macro for opencl

> error handling.

> 

> On 02/07/18 21:55, Ruiling Song wrote:

> > Signed-off-by: Ruiling Song <ruiling.song@intel.com>

> > ---

> >  libavfilter/opencl.h            |  4 ++--

> >  libavfilter/vf_avgblur_opencl.c | 45 +++++++++--------------------------

> >  libavfilter/vf_overlay_opencl.c | 29 +++++------------------

> >  libavfilter/vf_program_opencl.c | 14 ++---------

> >  libavfilter/vf_tonemap_opencl.c | 33 +++++---------------------

> >  libavfilter/vf_unsharp_opencl.c | 52 +++++++++--------------------------------

> >  6 files changed, 38 insertions(+), 139 deletions(-)

> 

> This doesn't apply - it seems like it's on top of another patch?

I am really sorry that I based on the patch on the old version and forget to merge them into one.
Have been fixed and sent to the mailing list again.

Thanks!
Ruiling
> 

> Code looks good to me.

> 

> Thanks,

> 

> - Mark

> 

>

Patch hide | download patch | download mbox

diff --git a/libavfilter/opencl.h b/libavfilter/opencl.h
index fd76f72..0ed360b 100644
--- a/libavfilter/opencl.h
+++ b/libavfilter/opencl.h
@@ -116,9 +116,9 @@  int ff_opencl_filter_work_size_from_image(AVFilterContext *avctx,
  * A helper macro to handle OpenCL error. It will assign errcode to
  * variable err, log error msg, and jump to fail label on error.
  */
-#define OCL_FAIL_ON_ERR(logctx, cle, errcode, ...) do {\
+#define CL_FAIL_ON_ERROR(errcode, ...) do {\
     if (cle != CL_SUCCESS) {\
-        av_log(logctx, AV_LOG_ERROR, __VA_ARGS__);\
+        av_log(avctx, AV_LOG_ERROR, __VA_ARGS__);\
         err = errcode;\
         goto fail;\
     }\
diff --git a/libavfilter/vf_avgblur_opencl.c b/libavfilter/vf_avgblur_opencl.c
index d1d3eb1..bc6bcab 100644
--- a/libavfilter/vf_avgblur_opencl.c
+++ b/libavfilter/vf_avgblur_opencl.c
@@ -64,26 +64,16 @@  static int avgblur_opencl_init(AVFilterContext *avctx)
     ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
                                               ctx->ocf.hwctx->device_id,
                                               0, &cle);
-    if (!ctx->command_queue) {
-        av_log(avctx, AV_LOG_ERROR, "Failed to create OpenCL "
-               "command queue: %d.\n", cle);
-        err = AVERROR(EIO);
-        goto fail;
-    }
+    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL "
+                     "command queue %d.\n", cle);
 
     ctx->kernel_horiz = clCreateKernel(ctx->ocf.program,"avgblur_horiz", &cle);
-    if (!ctx->kernel_horiz) {
-        av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle);
-        err = AVERROR(EIO);
-        goto fail;
-    }
+    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create horizontal "
+                     "kernel %d.\n", cle);
 
     ctx->kernel_vert = clCreateKernel(ctx->ocf.program,"avgblur_vert", &cle);
-    if (!ctx->kernel_vert) {
-        av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle);
-        err = AVERROR(EIO);
-        goto fail;
-    }
+    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create vertical "
+                     "kernel %d.\n", cle);
 
     ctx->initialised = 1;
     return 0;
@@ -236,12 +226,8 @@  static int avgblur_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
             cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel_horiz, 2, NULL,
                                          global_work, NULL,
                                          0, NULL, NULL);
-            if (cle != CL_SUCCESS) {
-                av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n",
-                       cle);
-                err = AVERROR(EIO);
-                goto fail;
-            }
+            CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue horizontal "
+                             "kernel: %d.\n", cle);
             cle = clFinish(ctx->command_queue);
 
             err = ff_opencl_filter_work_size_from_image(avctx, global_work,
@@ -259,22 +245,13 @@  static int avgblur_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
             cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel_vert, 2, NULL,
                                          global_work, NULL,
                                          0, NULL, NULL);
-            if (cle != CL_SUCCESS) {
-                av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n",
-                       cle);
-                err = AVERROR(EIO);
-                goto fail;
-            }
+            CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue vertical "
+                             "kernel: %d.\n", cle);
         }
     }
 
     cle = clFinish(ctx->command_queue);
-    if (cle != CL_SUCCESS) {
-        av_log(avctx, AV_LOG_ERROR, "Failed to finish command queue: %d.\n",
-               cle);
-        err = AVERROR(EIO);
-        goto fail;
-    }
+    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
 
     err = av_frame_copy_props(output, input);
     if (err < 0)
diff --git a/libavfilter/vf_overlay_opencl.c b/libavfilter/vf_overlay_opencl.c
index 556ce35..e9c8532 100644
--- a/libavfilter/vf_overlay_opencl.c
+++ b/libavfilter/vf_overlay_opencl.c
@@ -100,19 +100,11 @@  static int overlay_opencl_load(AVFilterContext *avctx,
     ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
                                               ctx->ocf.hwctx->device_id,
                                               0, &cle);
-    if (!ctx->command_queue) {
-        av_log(avctx, AV_LOG_ERROR, "Failed to create OpenCL "
-               "command queue: %d.\n", cle);
-        err = AVERROR(EIO);
-        goto fail;
-    }
+    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL "
+                     "command queue %d.\n", cle);
 
     ctx->kernel = clCreateKernel(ctx->ocf.program, kernel, &cle);
-    if (!ctx->kernel) {
-        av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle);
-        err = AVERROR(EIO);
-        goto fail;
-    }
+    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle);
 
     ctx->initialised = 1;
     return 0;
@@ -209,21 +201,12 @@  static int overlay_opencl_blend(FFFrameSync *fs)
 
         cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
                                      global_work, NULL, 0, NULL, NULL);
-        if (cle != CL_SUCCESS) {
-            av_log(avctx, AV_LOG_ERROR, "Failed to enqueue "
-                   "overlay kernel for plane %d: %d.\n", cle, plane);
-            err = AVERROR(EIO);
-            goto fail;
-        }
+        CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue overlay kernel "
+                         "for plane %d: %d.\n", plane, cle);
     }
 
     cle = clFinish(ctx->command_queue);
-    if (cle != CL_SUCCESS) {
-        av_log(avctx, AV_LOG_ERROR, "Failed to finish "
-               "command queue: %d.\n", cle);
-        err = AVERROR(EIO);
-        goto fail;
-    }
+    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
 
     err = av_frame_copy_props(output, input_main);
 
diff --git a/libavfilter/vf_program_opencl.c b/libavfilter/vf_program_opencl.c
index a002792..dfb2565 100644
--- a/libavfilter/vf_program_opencl.c
+++ b/libavfilter/vf_program_opencl.c
@@ -148,21 +148,11 @@  static int program_opencl_run(AVFilterContext *avctx)
 
         cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
                                      global_work, NULL, 0, NULL, NULL);
-        if (cle != CL_SUCCESS) {
-            av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n",
-                   cle);
-            err = AVERROR(EIO);
-            goto fail;
-        }
+        CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue kernel: %d.\n", cle);
     }
 
     cle = clFinish(ctx->command_queue);
-    if (cle != CL_SUCCESS) {
-        av_log(avctx, AV_LOG_ERROR, "Failed to finish command queue: %d.\n",
-               cle);
-        err = AVERROR(EIO);
-        goto fail;
-    }
+    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
 
     if (ctx->nb_inputs > 0) {
         err = av_frame_copy_props(output, ctx->frames[0]);
diff --git a/libavfilter/vf_tonemap_opencl.c b/libavfilter/vf_tonemap_opencl.c
index 36c7fbe..241f95e 100644
--- a/libavfilter/vf_tonemap_opencl.c
+++ b/libavfilter/vf_tonemap_opencl.c
@@ -262,29 +262,17 @@  static int tonemap_opencl_init(AVFilterContext *avctx)
     ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
                                               ctx->ocf.hwctx->device_id,
                                               0, &cle);
-    if (!ctx->command_queue) {
-        av_log(avctx, AV_LOG_ERROR, "Failed to create OpenCL "
-               "command queue: %d.\n", cle);
-        err = AVERROR(EIO);
-        goto fail;
-    }
+    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL "
+                     "command queue %d.\n", cle);
 
     ctx->kernel = clCreateKernel(ctx->ocf.program, "tonemap", &cle);
-    if (!ctx->kernel) {
-        av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle);
-        err = AVERROR(EIO);
-        goto fail;
-    }
+    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle);
 
     ctx->util_mem =
         clCreateBuffer(ctx->ocf.hwctx->context, 0,
                        (2 * DETECTION_FRAMES + 7) * sizeof(unsigned),
                        NULL, &cle);
-    if (cle != CL_SUCCESS) {
-        av_log(avctx, AV_LOG_ERROR, "Failed to create util buffer: %d.\n", cle);
-        err = AVERROR(EIO);
-        goto fail;
-    }
+    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create util buffer: %d.\n", cle);
 
     ctx->initialised = 1;
     return 0;
@@ -349,11 +337,7 @@  static int launch_kernel(AVFilterContext *avctx, cl_kernel kernel,
     cle = clEnqueueNDRangeKernel(ctx->command_queue, kernel, 2, NULL,
                                  global_work, local_work,
                                  0, NULL, NULL);
-    if (cle != CL_SUCCESS) {
-        av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n",
-               cle);
-        return AVERROR(EIO);
-    }
+    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue kernel: %d.\n", cle);
     return 0;
 fail:
     return err;
@@ -482,12 +466,7 @@  static int tonemap_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
     }
 
     cle = clFinish(ctx->command_queue);
-    if (cle != CL_SUCCESS) {
-        av_log(avctx, AV_LOG_ERROR, "Failed to finish command queue: %d.\n",
-               cle);
-        err = AVERROR(EIO);
-        goto fail;
-    }
+    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
 
     av_frame_free(&input);
 
diff --git a/libavfilter/vf_unsharp_opencl.c b/libavfilter/vf_unsharp_opencl.c
index 5b1eb59..d76d1b1 100644
--- a/libavfilter/vf_unsharp_opencl.c
+++ b/libavfilter/vf_unsharp_opencl.c
@@ -76,12 +76,8 @@  static int unsharp_opencl_init(AVFilterContext *avctx)
     ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
                                               ctx->ocf.hwctx->device_id,
                                               0, &cle);
-    if (!ctx->command_queue) {
-        av_log(avctx, AV_LOG_ERROR, "Failed to create OpenCL "
-               "command queue: %d.\n", cle);
-        err = AVERROR(EIO);
-        goto fail;
-    }
+    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL "
+                     "command queue %d.\n", cle);
 
     // Use global kernel if mask size will be too big for the local store..
     ctx->global = (ctx->luma_size_x   > 17.0f ||
@@ -92,11 +88,7 @@  static int unsharp_opencl_init(AVFilterContext *avctx)
     ctx->kernel = clCreateKernel(ctx->ocf.program,
                                  ctx->global ? "unsharp_global"
                                              : "unsharp_local", &cle);
-    if (!ctx->kernel) {
-        av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle);
-        err = AVERROR(EIO);
-        goto fail;
-    }
+    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle);
 
     ctx->initialised = 1;
     return 0;
@@ -176,12 +168,8 @@  static int unsharp_opencl_make_filter_params(AVFilterContext *avctx)
                                     CL_MEM_COPY_HOST_PTR |
                                     CL_MEM_HOST_NO_ACCESS,
                                     matrix_bytes, matrix, &cle);
-            if (!buffer) {
-                av_log(avctx, AV_LOG_ERROR, "Failed to create matrix buffer: "
-                       "%d.\n", cle);
-                err = AVERROR(EIO);
-                goto fail;
-            }
+            CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create matrix buffer: "
+                             "%d.\n", cle);
             ctx->plane[p].matrix = buffer;
         } else {
             buffer = clCreateBuffer(ctx->ocf.hwctx->context,
@@ -190,12 +178,8 @@  static int unsharp_opencl_make_filter_params(AVFilterContext *avctx)
                                     CL_MEM_HOST_NO_ACCESS,
                                     sizeof(ctx->plane[p].blur_x),
                                     ctx->plane[p].blur_x, &cle);
-            if (!buffer) {
-                av_log(avctx, AV_LOG_ERROR, "Failed to create x-coef buffer: "
-                       "%d.\n", cle);
-                err = AVERROR(EIO);
-                goto fail;
-            }
+            CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create x-coef buffer: "
+                             "%d.\n", cle);
             ctx->plane[p].coef_x = buffer;
 
             buffer = clCreateBuffer(ctx->ocf.hwctx->context,
@@ -204,12 +188,8 @@  static int unsharp_opencl_make_filter_params(AVFilterContext *avctx)
                                     CL_MEM_HOST_NO_ACCESS,
                                     sizeof(ctx->plane[p].blur_y),
                                     ctx->plane[p].blur_y, &cle);
-            if (!buffer) {
-                av_log(avctx, AV_LOG_ERROR, "Failed to create y-coef buffer: "
-                       "%d.\n", cle);
-                err = AVERROR(EIO);
-                goto fail;
-            }
+            CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create y-coef buffer: "
+                             "%d.\n", cle);
             ctx->plane[p].coef_y = buffer;
         }
 
@@ -296,21 +276,11 @@  static int unsharp_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
         cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
                                      global_work, ctx->global ? NULL : local_work,
                                      0, NULL, NULL);
-        if (cle != CL_SUCCESS) {
-            av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n",
-                   cle);
-            err = AVERROR(EIO);
-            goto fail;
-        }
+        CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue kernel: %d.\n", cle);
     }
 
     cle = clFinish(ctx->command_queue);
-    if (cle != CL_SUCCESS) {
-        av_log(avctx, AV_LOG_ERROR, "Failed to finish command queue: %d.\n",
-               cle);
-        err = AVERROR(EIO);
-        goto fail;
-    }
+    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
 
     err = av_frame_copy_props(output, input);
     if (err < 0)