[FFmpeg-devel] [PATCH] libavfilter/opencl.h: Add macro for setting opencl kernel
Danil Iashchenko
danyaschenko at gmail.com
Fri Jun 15 05:55:34 EEST 2018
---
Hi!
I like your idea with OCL_FAIL_ON_ERR(), but still do not know which one is better.
My idea relies on fact, that there are only few OpenCL functions which are used multiple times in filters:
clSetKernelArg, clCreateKernel(in case when there are multiple kernels) and maybe
clEnqueueNDRangeKernel. So that is why my purpose is totally wrap them and significantly reduce code,
but yes, there are some restrictions, like you can not use kernel_arg++ when setting kernel arguments.
And still most of cl-error checking statements appear after using cl-functions listed above.
Thanks, Danil
libavfilter/opencl.h | 15 ++++++++++
libavfilter/vf_convolution_opencl.c | 43 ++++------------------------
libavfilter/vf_overlay_opencl.c | 44 +++++++++++-----------------
libavfilter/vf_unsharp_opencl.c | 57 ++++++-------------------------------
4 files changed, 46 insertions(+), 113 deletions(-)
diff --git a/libavfilter/opencl.h b/libavfilter/opencl.h
index c0a4519..7441b11 100644
--- a/libavfilter/opencl.h
+++ b/libavfilter/opencl.h
@@ -46,6 +46,21 @@ typedef struct OpenCLFilterContext {
int output_height;
} OpenCLFilterContext;
+
+/**
+ * set argument to specific Kernel.
+ * This macro relies on usage of local label "fail" and variables:
+ * avctx, cle and err.
+ */
+#define CL_SET_KERNEL_ARG(kernel, arg_num, type, arg) \
+ cle = clSetKernelArg(kernel, arg_num, sizeof(type), arg); \
+ if (cle != CL_SUCCESS) { \
+ av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " \
+ "argument %d: error %d.\n", arg_num, cle); \
+ err = AVERROR(EIO); \
+ goto fail; \
+ }
+
/**
* Return that all inputs and outputs support only AV_PIX_FMT_OPENCL.
*/
diff --git a/libavfilter/vf_convolution_opencl.c b/libavfilter/vf_convolution_opencl.c
index 2df51e0..4d0ecf8 100644
--- a/libavfilter/vf_convolution_opencl.c
+++ b/libavfilter/vf_convolution_opencl.c
@@ -204,43 +204,12 @@ static int convolution_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->dims[p]);
- 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_mem), &ctx->matrix[p]);
- if (cle != CL_SUCCESS) {
- av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
- "matrix argument: %d.\n", cle);
- goto fail;
- }
- cle = clSetKernelArg(ctx->kernel, 4, sizeof(cl_float), &ctx->rdivs[p]);
- if (cle != CL_SUCCESS) {
- av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
- "rdiv argument: %d.\n", cle);
- goto fail;
- }
- cle = clSetKernelArg(ctx->kernel, 5, sizeof(cl_float), &ctx->biases[p]);
- if (cle != CL_SUCCESS) {
- av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
- "bias 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->dims[p]);
+ CL_SET_KERNEL_ARG(ctx->kernel, 3, cl_mem, &ctx->matrix[p]);
+ CL_SET_KERNEL_ARG(ctx->kernel, 4, cl_float, &ctx->rdivs[p]);
+ CL_SET_KERNEL_ARG(ctx->kernel, 5, cl_float, &ctx->biases[p]);
err = ff_opencl_filter_work_size_from_image(avctx, global_work, output, p, 0);
if (err < 0)
diff --git a/libavfilter/vf_overlay_opencl.c b/libavfilter/vf_overlay_opencl.c
index b43050d..556ce35 100644
--- a/libavfilter/vf_overlay_opencl.c
+++ b/libavfilter/vf_overlay_opencl.c
@@ -167,47 +167,39 @@ 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++;
}
err = ff_opencl_filter_work_size_from_image(avctx, global_work,
@@ -241,10 +233,6 @@ 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;
diff --git a/libavfilter/vf_unsharp_opencl.c b/libavfilter/vf_unsharp_opencl.c
index 19c9185..385d851 100644
--- a/libavfilter/vf_unsharp_opencl.c
+++ b/libavfilter/vf_unsharp_opencl.c
@@ -268,56 +268,17 @@ 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;
- }
+ 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) {
- 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;
- }
+ CL_SET_KERNEL_ARG(ctx->kernel, 5, cl_mem, &ctx->plane[p].matrix);
} 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, 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,
--
2.7.4
More information about the ffmpeg-devel
mailing list