float4 dstPix = convPix * div + bias;
write_imagef(dst, loc, dstPix);
}
+
+
+__kernel void sobel_global(__write_only image2d_t dst,
+ __read_only image2d_t src,
+ float div,
+ float bias)
+{
+ const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE |
+ CLK_ADDRESS_CLAMP_TO_EDGE |
+ CLK_FILTER_NEAREST);
+
+ int2 loc = (int2)(get_global_id(0), get_global_id(1));
+
+ float4 sum1 = read_imagef(src, sampler, loc + (int2)(-1,-1)) * -1 +
+ read_imagef(src, sampler, loc + (int2)( 0,-1)) * -2 +
+ read_imagef(src, sampler, loc + (int2)( 1,-1)) * -1 +
+ read_imagef(src, sampler, loc + (int2)(-1, 1)) * 1 +
+ read_imagef(src, sampler, loc + (int2)( 0, 1)) * 2 +
+ read_imagef(src, sampler, loc + (int2)( 1, 1)) * 1;
+
+ float4 sum2 = read_imagef(src, sampler, loc + (int2)(-1,-1)) * -1 +
+ read_imagef(src, sampler, loc + (int2)(-1, 0)) * -2 +
+ read_imagef(src, sampler, loc + (int2)(-1, 1)) * -1 +
+ read_imagef(src, sampler, loc + (int2)( 1,-1)) * 1 +
+ read_imagef(src, sampler, loc + (int2)( 1, 0)) * 2 +
+ read_imagef(src, sampler, loc + (int2)( 1, 1)) * 1;
+
+ float4 dstPix = hypot(sum1, sum2) * div + bias;
+ write_imagef(dst, loc, dstPix);
+}
+
+__kernel void prewitt_global(__write_only image2d_t dst,
+ __read_only image2d_t src,
+ float div,
+ float bias)
+{
+ const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE |
+ CLK_ADDRESS_CLAMP_TO_EDGE |
+ CLK_FILTER_NEAREST);
+
+ int2 loc = (int2)(get_global_id(0), get_global_id(1));
+
+ float4 sum1 = read_imagef(src, sampler, loc + (int2)(-1,-1)) * 1 +
+ read_imagef(src, sampler, loc + (int2)( 0,-1)) * 1 +
+ read_imagef(src, sampler, loc + (int2)( 1,-1)) * 1 +
+ read_imagef(src, sampler, loc + (int2)(-1, 1)) * -1 +
+ read_imagef(src, sampler, loc + (int2)( 0, 1)) * -1 +
+ read_imagef(src, sampler, loc + (int2)( 1, 1)) * -1;
+
+ float4 sum2 = read_imagef(src, sampler, loc + (int2)(-1,-1)) * 1 +
+ read_imagef(src, sampler, loc + (int2)(-1, 0)) * 1 +
+ read_imagef(src, sampler, loc + (int2)(-1, 1)) * 1 +
+ read_imagef(src, sampler, loc + (int2)( 1,-1)) * -1 +
+ read_imagef(src, sampler, loc + (int2)( 1, 0)) * -1 +
+ read_imagef(src, sampler, loc + (int2)( 1, 1)) * -1;
+
+ float4 dstPix = hypot(sum1, sum2) * div + bias;
+ write_imagef(dst, loc, dstPix);
+}
+
+__kernel void roberts_global(__write_only image2d_t dst,
+ __read_only image2d_t src,
+ float div,
+ float bias)
+{
+ const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE |
+ CLK_ADDRESS_CLAMP_TO_EDGE |
+ CLK_FILTER_NEAREST);
+
+ int2 loc = (int2)(get_global_id(0), get_global_id(1));
+
+ float4 sum1 = read_imagef(src, sampler, loc + (int2)(-1,-1)) * 1 +
+ read_imagef(src, sampler, loc + (int2)( 0,-1)) * -1;
+
+
+ float4 sum2 = read_imagef(src, sampler, loc + (int2)(-1, 0)) * -1 +
+ read_imagef(src, sampler, loc + (int2)( 0, 0)) * 1;
+
+
+ float4 dstPix = hypot(sum1, sum2) * div + bias;
+ write_imagef(dst, loc, dstPix);
+}
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;
goto fail;
}
- ctx->kernel = clCreateKernel(ctx->ocf.program, "convolution_global", &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);
if (!ctx->kernel) {
av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle);
err = AVERROR(EIO);
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),
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);
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);
+ if (cle != CL_SUCCESS) {
+ av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n",
+ cle);
+ err = AVERROR(EIO);
+ goto fail;
+ }
+ } 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);
+ if (cle != CL_SUCCESS) {
+ av_log(avctx, AV_LOG_ERROR, "Failed to copy plane %d: %d.\n",
+ p, cle);
+ err = AVERROR(EIO);
+ goto fail;
+ }
+ } 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);
+ if (cle != CL_SUCCESS) {
+ av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n",
+ cle);
+ err = AVERROR(EIO);
+ goto fail;
+ }
+ }
}
}
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 },
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,
+ .inputs = convolution_opencl_inputs,
+ .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,
.outputs = convolution_opencl_outputs,
.flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
};
+
+#endif /* CONFIG_ROBERTS_OPENCL_FILTER */