]> git.sesse.net Git - ffmpeg/blobdiff - libavfilter/vf_unsharp_opencl.c
projects / ffmpeg / blobdiff
? search:
summary | shortlog | log | commit | commitdiff | tree
raw | inline | side by side
avfilter: Constify all AVFilters
[ffmpeg] / libavfilter / vf_unsharp_opencl.c
diff --git a/libavfilter/vf_unsharp_opencl.c b/libavfilter/vf_unsharp_opencl.c
index 6a453c014bbe89ef61c818e33e649ca565cd61da..200350de9f2677ea65a9b5ca69137a88f22409be 100644 (file)
--- 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;
         }
 
@@ -268,68 +248,27 @@ static int unsharp_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
         if (!dst)
             break;
 
-        cle = clSetKernelArg(ctx->kernel, 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);
-            goto fail;
-        }
-        cle = clSetKernelArg(ctx->kernel, 1, sizeof(cl_mem), &src);
-        if (cle != CL_SUCCESS) {
-            av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
-                   "source image argument: %d.\n", cle);
-            goto fail;
-        }
-        cle = clSetKernelArg(ctx->kernel, 2, sizeof(cl_int), &ctx->plane[p].size_x);
-        if (cle != CL_SUCCESS) {
-            av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
-                   "matrix size argument: %d.\n", cle);
-            goto fail;
-        }
-        cle = clSetKernelArg(ctx->kernel, 3, sizeof(cl_int), &ctx->plane[p].size_y);
-        if (cle != CL_SUCCESS) {
-            av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
-                   "matrix size argument: %d.\n", cle);
-            goto fail;
-        }
-        cle = clSetKernelArg(ctx->kernel, 4, sizeof(cl_float), &ctx->plane[p].amount);
-        if (cle != CL_SUCCESS) {
-            av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
-                   "amount argument: %d.\n", cle);
-            goto fail;
-        }
-        if (ctx->global) {
-            cle = clSetKernelArg(ctx->kernel, 5, sizeof(cl_mem), &ctx->plane[p].matrix);
-            if (cle != CL_SUCCESS) {
-                av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
-                       "matrix argument: %d.\n", cle);
-                goto fail;
-            }
-        } else {
-            cle = clSetKernelArg(ctx->kernel, 5, sizeof(cl_mem), &ctx->plane[p].coef_x);
-            if (cle != CL_SUCCESS) {
-                av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
-                       "x-coef argument: %d.\n", cle);
-                goto fail;
-            }
-            cle = clSetKernelArg(ctx->kernel, 6, sizeof(cl_mem), &ctx->plane[p].coef_y);
-            if (cle != CL_SUCCESS) {
-                av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
-                       "y-coef argument: %d.\n", cle);
-                goto fail;
-            }
-        }
+        CL_SET_KERNEL_ARG(ctx->kernel, 0, cl_mem,   &dst);
+        CL_SET_KERNEL_ARG(ctx->kernel, 1, cl_mem,   &src);
+        CL_SET_KERNEL_ARG(ctx->kernel, 2, cl_int,   &ctx->plane[p].size_x);
+        CL_SET_KERNEL_ARG(ctx->kernel, 3, cl_int,   &ctx->plane[p].size_y);
+        CL_SET_KERNEL_ARG(ctx->kernel, 4, cl_float, &ctx->plane[p].amount);
 
         if (ctx->global) {
-            global_work[0] = output->width;
-            global_work[1] = output->height;
+            CL_SET_KERNEL_ARG(ctx->kernel, 5, cl_mem, &ctx->plane[p].matrix);
         } else {
-            global_work[0] = FFALIGN(output->width,  16);
-            global_work[1] = FFALIGN(output->height, 16);
-            local_work[0]  = 16;
-            local_work[1]  = 16;
+            CL_SET_KERNEL_ARG(ctx->kernel, 5, cl_mem, &ctx->plane[p].coef_x);
+            CL_SET_KERNEL_ARG(ctx->kernel, 6, cl_mem, &ctx->plane[p].coef_y);
         }
 
+        err = ff_opencl_filter_work_size_from_image(avctx, global_work, output, p,
+                                                    ctx->global ? 0 : 16);
+        if (err < 0)
+            goto fail;
+
+        local_work[0]  = 16;
+        local_work[1]  = 16;
+
         av_log(avctx, AV_LOG_DEBUG, "Run kernel on plane %d "
                "(%"SIZE_SPECIFIER"x%"SIZE_SPECIFIER").\n",
                p, global_work[0], global_work[1]);
@@ -337,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)
@@ -469,7 +398,7 @@ static const AVFilterPad unsharp_opencl_outputs[] = {
     { NULL }
 };
 
-AVFilter ff_vf_unsharp_opencl = {
+const AVFilter ff_vf_unsharp_opencl = {
     .name           = "unsharp_opencl",
     .description    = NULL_IF_CONFIG_SMALL("Apply unsharp mask to input video"),
     .priv_size      = sizeof(UnsharpOpenCLContext),
Unnamed repository; edit this file 'description' to name the repository.