WIP: working but very dirty.
Loads an arbitrary OpenCL kernel and runs it on the image.
---
libavfilter/Makefile | 1 +
libavfilter/allfilters.c | 1 +
libavfilter/vf_opencl_inplace.c | 280 ++++++++++++++++++++++++++++++++++++++++
3 files changed, 282 insertions(+)
create mode 100644 libavfilter/vf_opencl_inplace.c
diff --git a/libavfilter/Makefile b/libavfilter/Makefile
index 4ec244e..3bb6cc0 100644
--- a/libavfilter/Makefile
+++ b/libavfilter/Makefile
@@ -72,6 +72,7 @@ OBJS-$(CONFIG_NEGATE_FILTER) += vf_lut.o
OBJS-$(CONFIG_NOFORMAT_FILTER) += vf_format.o
OBJS-$(CONFIG_NULL_FILTER) += vf_null.o
OBJS-$(CONFIG_OCV_FILTER) += vf_libopencv.o
+OBJS-$(CONFIG_OPENCL_INPLACE_FILTER) += vf_opencl_inplace.o
OBJS-$(CONFIG_OVERLAY_FILTER) += vf_overlay.o
OBJS-$(CONFIG_PAD_FILTER) += vf_pad.o
OBJS-$(CONFIG_PIXDESCTEST_FILTER) += vf_pixdesctest.o
diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c
index 36a484c..1bbe1bd 100644
--- a/libavfilter/allfilters.c
+++ b/libavfilter/allfilters.c
@@ -95,6 +95,7 @@ void avfilter_register_all(void)
REGISTER_FILTER(NOFORMAT, noformat, vf);
REGISTER_FILTER(NULL, null, vf);
REGISTER_FILTER(OCV, ocv, vf);
+ REGISTER_FILTER(OPENCL_INPLACE, opencl_inplace, vf);
REGISTER_FILTER(OVERLAY, overlay, vf);
REGISTER_FILTER(PAD, pad, vf);
REGISTER_FILTER(PIXDESCTEST, pixdesctest, vf);
diff --git a/libavfilter/vf_opencl_inplace.c b/libavfilter/vf_opencl_inplace.c
new file mode 100644
index 0000000..9a15a39
--- /dev/null
+++ b/libavfilter/vf_opencl_inplace.c
@@ -0,0 +1,280 @@
+/*
+ * This file is part of Libav.
+ *
+ * Libav is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation; either
+ * version 2.1 of the License, or (at your option) any later version.
+ *
+ * Libav is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with Libav; if not, write to the Free Software
+ * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
+ */
+
+#include <stdio.h>
+#include <string.h>
+
+#include "libavutil/buffer.h"
+#include "libavutil/hwcontext.h"
+#include "libavutil/hwcontext_opencl.h"
+#include "libavutil/log.h"
+#include "libavutil/mem.h"
+#include "libavutil/pixdesc.h"
+#include "libavutil/opt.h"
+
+#include "avfilter.h"
+#include "formats.h"
+#include "internal.h"
+#include "video.h"
+
+typedef struct OpenCLInplaceContext {
+ const AVClass *class;
+
+ AVBufferRef *device_ref;
+ AVOpenCLDeviceContext *hwctx;
+
+ cl_command_queue command_queue;
+ cl_program program;
+ cl_kernel kernel;
+
+ char *program_file;
+ char *kernel_name;
+ int work_width;
+ int work_height;
+} OpenCLInplaceContext;
+
+static int opencl_inplace_query_formats(AVFilterContext *avctx)
+{
+ const static enum AVPixelFormat formats[] = {
+ AV_PIX_FMT_OPENCL,
+ AV_PIX_FMT_NONE,
+ };
+
+ ff_formats_ref(ff_make_format_list(formats),
+ &avctx->inputs[0]->out_formats);
+
+ ff_formats_ref(ff_make_format_list(formats),
+ &avctx->outputs[0]->in_formats);
+
+ return 0;
+}
+
+static int opencl_inplace_load_program(AVFilterContext *avctx)
+{
+ OpenCLInplaceContext *ctx = avctx->priv;
+ cl_int cle;
+ FILE *source_file;
+ char *source;
+ size_t source_length;
+ int err;
+
+ ctx->command_queue = clCreateCommandQueue(ctx->hwctx->context,
+ ctx->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(ENOSYS);
+ goto fail;
+ }
+
+ // How to get this in properly?
+ source_file = fopen(ctx->program_file, "r");
+ source = malloc(1 << 20);
+ source_length = fread(source, 1, 1 << 20, source_file);
+ fclose(source_file);
+
+ ctx->program = clCreateProgramWithSource(ctx->hwctx->context, 1,
+ (const char**)&source,
+ &source_length, &cle);
+ if (!ctx->program) {
+ av_log(avctx, AV_LOG_ERROR, "Failed to create program: %d.\n", cle);
+ err = AVERROR(EIO);
+ goto fail;
+ }
+
+ cle = clBuildProgram(ctx->program, 1, &ctx->hwctx->device_id,
+ NULL, NULL, NULL);
+ if (cle != CL_SUCCESS) {
+ av_log(avctx, AV_LOG_ERROR, "Failed to build program: %d.\n", cle);
+
+ if (cle == CL_BUILD_PROGRAM_FAILURE) {
+ char *log;
+ size_t log_length;
+
+ clGetProgramBuildInfo(ctx->program, ctx->hwctx->device_id,
+ CL_PROGRAM_BUILD_LOG, 0, NULL, &log_length);
+
+ log = av_malloc(log_length);
+ if (log) {
+ cle = clGetProgramBuildInfo(ctx->program,
+ ctx->hwctx->device_id,
+ CL_PROGRAM_BUILD_LOG,
+ log_length, log, NULL);
+ if (cle == CL_SUCCESS)
+ av_log(avctx, AV_LOG_ERROR, "Build log:\n%s\n", log);
+ }
+
+ av_free(log);
+ }
+
+ err = AVERROR(EIO);
+ goto fail;
+ }
+
+ ctx->kernel = clCreateKernel(ctx->program, ctx->kernel_name, &cle);
+ if (!ctx->kernel) {
+ av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle);
+ err = AVERROR(EIO);
+ goto fail;
+ }
+
+ return 0;
+
+fail:
+ return err;
+}
+
+static int opencl_inplace_filter_frame(AVFilterLink *inlink, AVFrame *pic)
+{
+ AVFilterContext *avctx = inlink->dst;
+ AVFilterLink *outlink = avctx->outputs[0];
+ OpenCLInplaceContext *ctx = avctx->priv;
+ AVHWFramesContext *hwfc;
+ cl_mem luma;
+ cl_int cle;
+ size_t global_work[2];
+ size_t local_work[2];
+ int err;
+
+ // Would be better if this could be done at init-time.
+ if (!pic->hw_frames_ctx)
+ return AVERROR(EINVAL);
+ hwfc = (AVHWFramesContext*)pic->hw_frames_ctx->data;
+ if (!hwfc || hwfc->format != AV_PIX_FMT_OPENCL)
+ return AVERROR(EINVAL);
+ ctx->hwctx = hwfc->device_ctx->hwctx;
+
+ if (!ctx->kernel) {
+ err = opencl_inplace_load_program(avctx);
+ if (err)
+ goto fail;
+ }
+
+ luma = (cl_mem)pic->data[0];
+ av_log(avctx, AV_LOG_DEBUG, "Luma input is %p.\n", luma);
+
+ cle = clSetKernelArg(ctx->kernel, 0, sizeof(cl_mem), &luma);
+ if (cle != CL_SUCCESS) {
+ av_log(avctx, AV_LOG_ERROR, "Failed to set kernel argument: %d.\n",
cle);
+ err = AVERROR(EIO);
+ goto fail;
+ }
+
+ global_work[0] = pic->width;
+ global_work[1] = pic->height;
+ local_work[0] = ctx->work_width;
+ local_work[1] = ctx->work_height;
+
+ av_log(avctx, AV_LOG_DEBUG, "Global %zux%zu, local %zux%zu.\n",
+ global_work[0], global_work[1],
+ local_work[0], local_work[1]);
+
+ cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
+ global_work, local_work, 0, NULL, NULL);
+ if (cle < 0) {
+ av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n",
+ cle);
+ err = AVERROR(EIO);
+ goto fail;
+ }
+
+ 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;
+ }
+
+ return ff_filter_frame(outlink, pic);
+
+fail:
+ // Missing cleanup.
+ return err;
+}
+
+static av_cold void opencl_inplace_uninit(AVFilterContext *avctx)
+{
+ OpenCLInplaceContext *ctx = avctx->priv;
+ cl_int cle;
+
+ cle = clReleaseKernel(ctx->kernel);
+ if (cle != CL_SUCCESS)
+ av_log(avctx, AV_LOG_ERROR, "Failed to release kernel: %d.\n",
+ cle);
+
+ cle = clReleaseProgram(ctx->program);
+ if (cle != CL_SUCCESS)
+ av_log(avctx, AV_LOG_ERROR, "Failed to release program: %d.\n",
+ cle);
+
+ cle = clReleaseCommandQueue(ctx->command_queue);
+ if (cle != CL_SUCCESS)
+ av_log(avctx, AV_LOG_ERROR, "Failed to release command queue: %d.\n",
+ cle);
+}
+
+#define OFFSET(x) offsetof(OpenCLInplaceContext, x)
+#define FLAGS (AV_OPT_FLAG_VIDEO_PARAM)
+static const AVOption opencl_inplace_options[] = {
+ { "program", "OpenCL program source file",
+ OFFSET(program_file), AV_OPT_TYPE_STRING, { .str = "test.cl" }, .flags =
FLAGS },
+ { "kernel", "Kernel name",
+ OFFSET(kernel_name), AV_OPT_TYPE_STRING, { .str = "test" }, .flags =
FLAGS },
+ { "work_width", "Width of a kernel work element",
+ OFFSET(work_width), AV_OPT_TYPE_INT, { .i64 = 0 }, 0, INT_MAX, .flags =
FLAGS },
+ { "work_height", "Height of a kernel work element",
+ OFFSET(work_height), AV_OPT_TYPE_INT, { .i64 = 0 }, 0, INT_MAX, .flags =
FLAGS },
+ { NULL },
+};
+
+static const AVClass opencl_inplace_class = {
+ .class_name = "opencl_inplace",
+ .item_name = av_default_item_name,
+ .option = opencl_inplace_options,
+ .version = LIBAVUTIL_VERSION_INT,
+};
+
+static const AVFilterPad opencl_inplace_inputs[] = {
+ {
+ .name = "default",
+ .type = AVMEDIA_TYPE_VIDEO,
+ .filter_frame = &opencl_inplace_filter_frame,
+ },
+ { NULL }
+};
+
+static const AVFilterPad opencl_inplace_outputs[] = {
+ {
+ .name = "default",
+ .type = AVMEDIA_TYPE_VIDEO,
+ },
+ { NULL }
+};
+
+AVFilter ff_vf_opencl_inplace = {
+ .name = "opencl_inplace",
+ .description = NULL_IF_CONFIG_SMALL("Do something unclear to an OpenCL
frame"),
+ .priv_size = sizeof(OpenCLInplaceContext),
+ .priv_class = &opencl_inplace_class,
+ .uninit = &opencl_inplace_uninit,
+ .query_formats = &opencl_inplace_query_formats,
+ .inputs = opencl_inplace_inputs,
+ .outputs = opencl_inplace_outputs,
+};
--
2.8.1
_______________________________________________
libav-devel mailing list
libav-devel@libav.org
https://lists.libav.org/mailman/listinfo/libav-devel
