X-Git-Url: https://git.sesse.net/?a=blobdiff_plain;f=libavfilter%2Fvf_convolution_opencl.c;h=00246b2e435f66136eb092809cc8dd62c4e8ad75;hb=23ce57af3ad684363881cdd66d5724f40963a65c;hp=4d0ecf8b9ae53e12d76e0524892cd40d189e3b74;hpb=b41b6b323417ce1e79bf818f91d60526697954f1;p=ffmpeg diff --git a/libavfilter/vf_convolution_opencl.c b/libavfilter/vf_convolution_opencl.c index 4d0ecf8b9ae..00246b2e435 100644 --- a/libavfilter/vf_convolution_opencl.c +++ b/libavfilter/vf_convolution_opencl.c @@ -47,12 +47,16 @@ typedef struct ConvolutionOpenCLContext { cl_float rdivs[4]; cl_float biases[4]; -} ConvolutionOpenCLContext; + cl_int planes; + cl_float scale; + cl_float delta; +} ConvolutionOpenCLContext; static int convolution_opencl_init(AVFilterContext *avctx) { ConvolutionOpenCLContext *ctx = avctx->priv; + const char *kernel_name; cl_int cle; int err; @@ -63,19 +67,21 @@ static int convolution_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; - } - - ctx->kernel = clCreateKernel(ctx->ocf.program, "convolution_global", &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 OpenCL " + "command queue %d.\n", cle); + + if (!strcmp(avctx->filter->name, "convolution_opencl")) { + kernel_name = "convolution_global"; + } else if (!strcmp(avctx->filter->name, "sobel_opencl")) { + kernel_name = "sobel_global"; + } else if (!strcmp(avctx->filter->name, "prewitt_opencl")){ + kernel_name = "prewitt_global"; + } else if (!strcmp(avctx->filter->name, "roberts_opencl")){ + kernel_name = "roberts_global"; } + ctx->kernel = clCreateKernel(ctx->ocf.program, kernel_name, &cle); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create " + "kernel %d.\n", cle); ctx->initialised = 1; return 0; @@ -173,6 +179,8 @@ static int convolution_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input) size_t global_work[2]; cl_mem src, dst; int err, p; + size_t origin[3] = {0, 0, 0}; + size_t region[3] = {0, 0, 1}; av_log(ctx, AV_LOG_DEBUG, "Filter input: %s, %ux%u (%"PRId64").\n", av_get_pix_fmt_name(input->format), @@ -186,9 +194,14 @@ static int convolution_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input) if (err < 0) goto fail; - err = convolution_opencl_make_filter_params(avctx); - if (err < 0) - goto fail; + if (!strcmp(avctx->filter->name, "convolution_opencl")) { + err = convolution_opencl_make_filter_params(avctx); + if (err < 0) + goto fail; + } else { + ctx->delta /= 255.0; + } + } output = ff_get_video_buffer(outlink, outlink->w, outlink->h); @@ -204,39 +217,62 @@ static int convolution_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input) if (!dst) break; - 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) - goto fail; - - av_log(avctx, AV_LOG_DEBUG, "Run kernel on plane %d " - "(%"SIZE_SPECIFIER"x%"SIZE_SPECIFIER").\n", - p, global_work[0], global_work[1]); - - 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; + if (!strcmp(avctx->filter->name, "convolution_opencl")) { + 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) + goto fail; + + av_log(avctx, AV_LOG_DEBUG, "Run kernel on plane %d " + "(%"SIZE_SPECIFIER"x%"SIZE_SPECIFIER").\n", + p, global_work[0], global_work[1]); + + cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL, + global_work, NULL, + 0, NULL, NULL); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue " + "kernel: %d.\n", cle); + } else { + if (!(ctx->planes & (1 << p))) { + err = ff_opencl_filter_work_size_from_image(avctx, region, output, p, 0); + if (err < 0) + goto fail; + + cle = clEnqueueCopyImage(ctx->command_queue, src, dst, + origin, origin, region, 0, NULL, NULL); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to copy plane %d: %d.\n", + p, cle); + } else { + 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_float, &ctx->scale); + CL_SET_KERNEL_ARG(ctx->kernel, 3, cl_float, &ctx->delta); + + err = ff_opencl_filter_work_size_from_image(avctx, global_work, output, p, 0); + if (err < 0) + goto fail; + + av_log(avctx, AV_LOG_DEBUG, "Run kernel on plane %d " + "(%"SIZE_SPECIFIER"x%"SIZE_SPECIFIER").\n", + p, global_work[0], global_work[1]); + + cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL, + global_work, NULL, + 0, NULL, NULL); + 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) @@ -284,8 +320,30 @@ static av_cold void convolution_opencl_uninit(AVFilterContext *avctx) ff_opencl_filter_uninit(avctx); } +static const AVFilterPad convolution_opencl_inputs[] = { + { + .name = "default", + .type = AVMEDIA_TYPE_VIDEO, + .filter_frame = &convolution_opencl_filter_frame, + .config_props = &ff_opencl_filter_config_input, + }, + { NULL } +}; + +static const AVFilterPad convolution_opencl_outputs[] = { + { + .name = "default", + .type = AVMEDIA_TYPE_VIDEO, + .config_props = &ff_opencl_filter_config_output, + }, + { NULL } +}; + #define OFFSET(x) offsetof(ConvolutionOpenCLContext, x) #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM) + +#if CONFIG_CONVOLUTION_OPENCL_FILTER + static const AVOption convolution_opencl_options[] = { { "0m", "set matrix for 2nd plane", OFFSET(matrix_str[0]), AV_OPT_TYPE_STRING, {.str="0 0 0 0 1 0 0 0 0"}, 0, 0, FLAGS }, { "1m", "set matrix for 2nd plane", OFFSET(matrix_str[1]), AV_OPT_TYPE_STRING, {.str="0 0 0 0 1 0 0 0 0"}, 0, 0, FLAGS }, @@ -304,30 +362,63 @@ static const AVOption convolution_opencl_options[] = { AVFILTER_DEFINE_CLASS(convolution_opencl); -static const AVFilterPad convolution_opencl_inputs[] = { - { - .name = "default", - .type = AVMEDIA_TYPE_VIDEO, - .filter_frame = &convolution_opencl_filter_frame, - .config_props = &ff_opencl_filter_config_input, - }, +AVFilter ff_vf_convolution_opencl = { + .name = "convolution_opencl", + .description = NULL_IF_CONFIG_SMALL("Apply convolution mask to input video"), + .priv_size = sizeof(ConvolutionOpenCLContext), + .priv_class = &convolution_opencl_class, + .init = &ff_opencl_filter_init, + .uninit = &convolution_opencl_uninit, + .query_formats = &ff_opencl_filter_query_formats, + .inputs = convolution_opencl_inputs, + .outputs = convolution_opencl_outputs, + .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE, +}; + +#endif /* CONFIG_CONVOLUTION_OPENCL_FILTER */ + +#if CONFIG_SOBEL_OPENCL_FILTER + +static const AVOption sobel_opencl_options[] = { + { "planes", "set planes to filter", OFFSET(planes), AV_OPT_TYPE_INT, {.i64=15}, 0, 15, FLAGS}, + { "scale", "set scale", OFFSET(scale), AV_OPT_TYPE_FLOAT, {.dbl=1.0}, 0.0, 65535, FLAGS}, + { "delta", "set delta", OFFSET(delta), AV_OPT_TYPE_FLOAT, {.dbl=0}, -65535, 65535, FLAGS}, { NULL } }; -static const AVFilterPad convolution_opencl_outputs[] = { - { - .name = "default", - .type = AVMEDIA_TYPE_VIDEO, - .config_props = &ff_opencl_filter_config_output, - }, +AVFILTER_DEFINE_CLASS(sobel_opencl); + +AVFilter ff_vf_sobel_opencl = { + .name = "sobel_opencl", + .description = NULL_IF_CONFIG_SMALL("Apply sobel operator"), + .priv_size = sizeof(ConvolutionOpenCLContext), + .priv_class = &sobel_opencl_class, + .init = &ff_opencl_filter_init, + .uninit = &convolution_opencl_uninit, + .query_formats = &ff_opencl_filter_query_formats, + .inputs = convolution_opencl_inputs, + .outputs = convolution_opencl_outputs, + .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE, +}; + +#endif /* CONFIG_SOBEL_OPENCL_FILTER */ + +#if CONFIG_PREWITT_OPENCL_FILTER + +static const AVOption prewitt_opencl_options[] = { + { "planes", "set planes to filter", OFFSET(planes), AV_OPT_TYPE_INT, {.i64=15}, 0, 15, FLAGS}, + { "scale", "set scale", OFFSET(scale), AV_OPT_TYPE_FLOAT, {.dbl=1.0}, 0.0, 65535, FLAGS}, + { "delta", "set delta", OFFSET(delta), AV_OPT_TYPE_FLOAT, {.dbl=0}, -65535, 65535, FLAGS}, { NULL } }; -AVFilter ff_vf_convolution_opencl = { - .name = "convolution_opencl", - .description = NULL_IF_CONFIG_SMALL("Apply convolution mask to input video"), +AVFILTER_DEFINE_CLASS(prewitt_opencl); + +AVFilter ff_vf_prewitt_opencl = { + .name = "prewitt_opencl", + .description = NULL_IF_CONFIG_SMALL("Apply prewitt operator"), .priv_size = sizeof(ConvolutionOpenCLContext), - .priv_class = &convolution_opencl_class, + .priv_class = &prewitt_opencl_class, .init = &ff_opencl_filter_init, .uninit = &convolution_opencl_uninit, .query_formats = &ff_opencl_filter_query_formats, @@ -335,3 +426,31 @@ AVFilter ff_vf_convolution_opencl = { .outputs = convolution_opencl_outputs, .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE, }; + +#endif /* CONFIG_PREWITT_OPENCL_FILTER */ + +#if CONFIG_ROBERTS_OPENCL_FILTER + +static const AVOption roberts_opencl_options[] = { + { "planes", "set planes to filter", OFFSET(planes), AV_OPT_TYPE_INT, {.i64=15}, 0, 15, FLAGS}, + { "scale", "set scale", OFFSET(scale), AV_OPT_TYPE_FLOAT, {.dbl=1.0}, 0.0, 65535, FLAGS}, + { "delta", "set delta", OFFSET(delta), AV_OPT_TYPE_FLOAT, {.dbl=0}, -65535, 65535, FLAGS}, + { NULL } +}; + +AVFILTER_DEFINE_CLASS(roberts_opencl); + +AVFilter ff_vf_roberts_opencl = { + .name = "roberts_opencl", + .description = NULL_IF_CONFIG_SMALL("Apply roberts operator"), + .priv_size = sizeof(ConvolutionOpenCLContext), + .priv_class = &roberts_opencl_class, + .init = &ff_opencl_filter_init, + .uninit = &convolution_opencl_uninit, + .query_formats = &ff_opencl_filter_query_formats, + .inputs = convolution_opencl_inputs, + .outputs = convolution_opencl_outputs, + .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE, +}; + +#endif /* CONFIG_ROBERTS_OPENCL_FILTER */