X-Git-Url: https://git.sesse.net/?a=blobdiff_plain;f=libavfilter%2Fvf_overlay_opencl.c;h=e9c853203b4a00de2bea15348068bd3f39b08453;hb=d7e0d428faaa04e2fd850eca82f314ca2ad3dfe5;hp=46ce42df847b3e391bc36b4e3e868073c1f56dea;hpb=5fc649505b2ddb72eee3fcb8611d3d715cd0f718;p=ffmpeg

diff --git a/libavfilter/vf_overlay_opencl.c b/libavfilter/vf_overlay_opencl.c
index 46ce42df847..e9c853203b4 100644
--- a/libavfilter/vf_overlay_opencl.c
+++ b/libavfilter/vf_overlay_opencl.c
@@ -16,16 +16,10 @@
  * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
  */
 
-#include "libavutil/avassert.h"
-#include "libavutil/buffer.h"
-#include "libavutil/common.h"
-#include "libavutil/hwcontext.h"
-#include "libavutil/hwcontext_opencl.h"
 #include "libavutil/log.h"
-#include "libavutil/mathematics.h"
 #include "libavutil/mem.h"
-#include "libavutil/pixdesc.h"
 #include "libavutil/opt.h"
+#include "libavutil/pixdesc.h"
 
 #include "avfilter.h"
 #include "framesync.h"
@@ -106,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;
@@ -173,69 +159,54 @@ static int overlay_opencl_blend(FFFrameSync *fs)
         kernel_arg = 0;
 
         mem = (cl_mem)output->data[plane];
-        cle = clSetKernelArg(ctx->kernel, kernel_arg++, sizeof(cl_mem), &mem);
-        if (cle != CL_SUCCESS)
-            goto fail_kernel_arg;
+        CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
+        kernel_arg++;
 
         mem = (cl_mem)input_main->data[plane];
-        cle = clSetKernelArg(ctx->kernel, kernel_arg++, sizeof(cl_mem), &mem);
-        if (cle != CL_SUCCESS)
-            goto fail_kernel_arg;
+        CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
+        kernel_arg++;
 
         mem = (cl_mem)input_overlay->data[plane];
-        cle = clSetKernelArg(ctx->kernel, kernel_arg++, sizeof(cl_mem), &mem);
-        if (cle != CL_SUCCESS)
-            goto fail_kernel_arg;
+        CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
+        kernel_arg++;
 
         if (ctx->alpha_separate) {
             mem = (cl_mem)input_overlay->data[ctx->nb_planes];
-            cle = clSetKernelArg(ctx->kernel, kernel_arg++, sizeof(cl_mem), &mem);
-            if (cle != CL_SUCCESS)
-                goto fail_kernel_arg;
+            CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
+            kernel_arg++;
         }
 
         x = ctx->x_position / (plane == 0 ? 1 : ctx->x_subsample);
         y = ctx->y_position / (plane == 0 ? 1 : ctx->y_subsample);
 
-        cle = clSetKernelArg(ctx->kernel, kernel_arg++, sizeof(cl_int), &x);
-        if (cle != CL_SUCCESS)
-            goto fail_kernel_arg;
-        cle = clSetKernelArg(ctx->kernel, kernel_arg++, sizeof(cl_int), &y);
-        if (cle != CL_SUCCESS)
-            goto fail_kernel_arg;
+        CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_int, &x);
+        kernel_arg++;
+        CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_int, &y);
+        kernel_arg++;
 
         if (ctx->alpha_separate) {
             cl_int alpha_adj_x = plane == 0 ? 1 : ctx->x_subsample;
             cl_int alpha_adj_y = plane == 0 ? 1 : ctx->y_subsample;
 
-            cle = clSetKernelArg(ctx->kernel, kernel_arg++, sizeof(cl_int), &alpha_adj_x);
-            if (cle != CL_SUCCESS)
-                goto fail_kernel_arg;
-            cle = clSetKernelArg(ctx->kernel, kernel_arg++, sizeof(cl_int), &alpha_adj_y);
-            if (cle != CL_SUCCESS)
-                goto fail_kernel_arg;
+            CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_int, &alpha_adj_x);
+            kernel_arg++;
+            CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_int, &alpha_adj_y);
+            kernel_arg++;
         }
 
-        global_work[0] = output->width;
-        global_work[1] = output->height;
+        err = ff_opencl_filter_work_size_from_image(avctx, global_work,
+                                                    output, plane, 0);
+        if (err < 0)
+            goto fail;
 
         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);
 
@@ -245,11 +216,8 @@ static int overlay_opencl_blend(FFFrameSync *fs)
 
     return ff_filter_frame(outlink, output);
 
-fail_kernel_arg:
-    av_log(avctx, AV_LOG_ERROR, "Failed to set kernel arg %d: %d.\n",
-           kernel_arg, cle);
-    err = AVERROR(EIO);
 fail:
+    av_frame_free(&output);
     return err;
 }