|
|
@@ -0,0 +1,443 @@
|
|
|
+/**
|
|
|
+ * SECTION:ufo-filter-task
|
|
|
+ * @Short_description: Process arbitrary Filter kernels
|
|
|
+ * @Title: filter
|
|
|
+ *
|
|
|
+ * This module is used to load an arbitrary #UfoPadding2DTask:kernel from
|
|
|
+ * #UfoPadding2DTask:filename and execute it on each input. The kernel must have
|
|
|
+ * only two global float array parameters, the first represents the input, the
|
|
|
+ * second one the output. #UfoPadding2DTask:num-dims must be changed, if the kernel
|
|
|
+ * accesses either one or three dimensional index spaces.
|
|
|
+ */
|
|
|
+
|
|
|
+#ifdef __APPLE__
|
|
|
+#include <OpenCL/cl.h>
|
|
|
+#else
|
|
|
+#include <CL/cl.h>
|
|
|
+#endif
|
|
|
+#include <ufo-gpu-task-iface.h>
|
|
|
+#include "ufo-padding-2d-task.h"
|
|
|
+
|
|
|
+typedef enum {
|
|
|
+ PADDING_ZERO = 0,
|
|
|
+ PADDING_CONST,
|
|
|
+ PADDING_GAVG,
|
|
|
+ PADDING_BREP
|
|
|
+} PaddingMode;
|
|
|
+
|
|
|
+struct _UfoPadding2DTaskPrivate {
|
|
|
+ guint in_width;
|
|
|
+ guint in_height;
|
|
|
+ guint out_width;
|
|
|
+ guint out_height;
|
|
|
+
|
|
|
+ // extent adds
|
|
|
+ guint xl;
|
|
|
+ guint xr;
|
|
|
+ guint yt;
|
|
|
+ guint yb;
|
|
|
+
|
|
|
+ size_t global_work_size_small[2];
|
|
|
+ size_t global_work_size_large[2];
|
|
|
+
|
|
|
+ PaddingMode mode;
|
|
|
+ // padding constant
|
|
|
+ float pconst;
|
|
|
+ cl_kernel kernel_iconst;
|
|
|
+ cl_kernel kernel_cpyimg;
|
|
|
+ cl_kernel kernel_brep;
|
|
|
+};
|
|
|
+
|
|
|
+static void ufo_task_interface_init (UfoTaskIface *iface);
|
|
|
+static void ufo_gpu_task_interface_init (UfoGpuTaskIface *iface);
|
|
|
+
|
|
|
+G_DEFINE_TYPE_WITH_CODE (UfoPadding2DTask, ufo_padding_2d_task, UFO_TYPE_TASK_NODE,
|
|
|
+ G_IMPLEMENT_INTERFACE (UFO_TYPE_TASK,
|
|
|
+ ufo_task_interface_init)
|
|
|
+ G_IMPLEMENT_INTERFACE (UFO_TYPE_GPU_TASK,
|
|
|
+ ufo_gpu_task_interface_init))
|
|
|
+
|
|
|
+#define UFO_PADDING_2D_TASK_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_PADDING_2D_TASK, UfoPadding2DTaskPrivate))
|
|
|
+
|
|
|
+enum {
|
|
|
+ PROP_0,
|
|
|
+ PROP_XL,
|
|
|
+ PROP_XR,
|
|
|
+ PROP_YT,
|
|
|
+ PROP_YB,
|
|
|
+ PROP_MODE,
|
|
|
+ PROP_PCONST,
|
|
|
+ N_PROPERTIES
|
|
|
+};
|
|
|
+
|
|
|
+static GParamSpec *properties[N_PROPERTIES] = { NULL, };
|
|
|
+
|
|
|
+UfoNode *
|
|
|
+ufo_padding_2d_task_new (void)
|
|
|
+{
|
|
|
+ return UFO_NODE (g_object_new (UFO_TYPE_PADDING_2D_TASK, NULL));
|
|
|
+}
|
|
|
+
|
|
|
+static gboolean
|
|
|
+ufo_padding_2d_task_process (UfoGpuTask *task,
|
|
|
+ UfoBuffer **inputs,
|
|
|
+ UfoBuffer *output,
|
|
|
+ UfoRequisition *requisition,
|
|
|
+ UfoGpuNode *node)
|
|
|
+{
|
|
|
+ UfoPadding2DTaskPrivate *priv;
|
|
|
+ cl_command_queue cmd_queue;
|
|
|
+ cl_mem in_mem;
|
|
|
+ cl_mem out_mem;
|
|
|
+
|
|
|
+ priv = UFO_PADDING_2D_TASK (task)->priv;
|
|
|
+
|
|
|
+ const PaddingMode mode = priv->mode;
|
|
|
+ const guint pxl = priv->xl;
|
|
|
+ const guint pyt = priv->yt;
|
|
|
+ float pval = priv->pconst;
|
|
|
+
|
|
|
+ const guint ixs = priv->in_width;
|
|
|
+ const guint iys = priv->in_height;
|
|
|
+ const guint oxs = priv->out_width;
|
|
|
+
|
|
|
+ if (mode == PADDING_GAVG) {
|
|
|
+ gfloat *indata = ufo_buffer_get_host_array (inputs[0], NULL);
|
|
|
+ gfloat sum = 0;
|
|
|
+ guint psz = ixs * iys;
|
|
|
+
|
|
|
+ for (guint i =0; i < psz; i++)
|
|
|
+ sum += indata[i];
|
|
|
+
|
|
|
+ pval = sum / (gfloat) psz;
|
|
|
+ }
|
|
|
+
|
|
|
+ cmd_queue = ufo_gpu_node_get_cmd_queue (node);
|
|
|
+ in_mem = ufo_buffer_get_device_array (inputs[0], cmd_queue);
|
|
|
+ out_mem = ufo_buffer_get_device_array (output, cmd_queue);
|
|
|
+
|
|
|
+ if ((mode == PADDING_ZERO) || (mode == PADDING_CONST) || (mode == PADDING_GAVG)) {
|
|
|
+ cl_kernel k_iconst = priv->kernel_iconst;
|
|
|
+ cl_kernel k_cpyimg = priv->kernel_cpyimg;
|
|
|
+
|
|
|
+ /// fill with constant
|
|
|
+ UFO_RESOURCES_CHECK_CLERR(clSetKernelArg( k_iconst, 0, sizeof(cl_mem), (void *) &out_mem));
|
|
|
+ UFO_RESOURCES_CHECK_CLERR(clSetKernelArg( k_iconst, 1, sizeof(int), &oxs));
|
|
|
+ UFO_RESOURCES_CHECK_CLERR(clSetKernelArg( k_iconst, 2, sizeof(float), &pval));
|
|
|
+
|
|
|
+ UFO_RESOURCES_CHECK_CLERR(clEnqueueNDRangeKernel(cmd_queue, k_iconst,
|
|
|
+ 2, NULL, priv->global_work_size_large, NULL,
|
|
|
+ 0, NULL, NULL));
|
|
|
+
|
|
|
+ /// copy old image
|
|
|
+ UFO_RESOURCES_CHECK_CLERR(clSetKernelArg( k_cpyimg, 0, sizeof(cl_mem), (void *) &in_mem));
|
|
|
+ UFO_RESOURCES_CHECK_CLERR(clSetKernelArg( k_cpyimg, 1, sizeof(cl_mem), (void *) &out_mem));
|
|
|
+ UFO_RESOURCES_CHECK_CLERR(clSetKernelArg( k_cpyimg, 2, sizeof(int), &ixs));
|
|
|
+ UFO_RESOURCES_CHECK_CLERR(clSetKernelArg( k_cpyimg, 3, sizeof(int), &oxs));
|
|
|
+ UFO_RESOURCES_CHECK_CLERR(clSetKernelArg( k_cpyimg, 4, sizeof(int), &pxl));
|
|
|
+ UFO_RESOURCES_CHECK_CLERR(clSetKernelArg( k_cpyimg, 5, sizeof(int), &pyt));
|
|
|
+
|
|
|
+ UFO_RESOURCES_CHECK_CLERR(clEnqueueNDRangeKernel(cmd_queue, k_cpyimg,
|
|
|
+ 2, NULL, priv->global_work_size_small, NULL,
|
|
|
+ 0, NULL, NULL));
|
|
|
+ }
|
|
|
+
|
|
|
+ if (mode == PADDING_BREP) {
|
|
|
+ cl_kernel k_brep = priv->kernel_brep;
|
|
|
+
|
|
|
+ UFO_RESOURCES_CHECK_CLERR(clSetKernelArg( k_brep, 0, sizeof(cl_mem), (void *) &in_mem));
|
|
|
+ UFO_RESOURCES_CHECK_CLERR(clSetKernelArg( k_brep, 1, sizeof(cl_mem), (void *) &out_mem));
|
|
|
+ UFO_RESOURCES_CHECK_CLERR(clSetKernelArg( k_brep, 2, sizeof(int), &ixs));
|
|
|
+ UFO_RESOURCES_CHECK_CLERR(clSetKernelArg( k_brep, 3, sizeof(int), &iys));
|
|
|
+ UFO_RESOURCES_CHECK_CLERR(clSetKernelArg( k_brep, 4, sizeof(int), &oxs));
|
|
|
+ UFO_RESOURCES_CHECK_CLERR(clSetKernelArg( k_brep, 5, sizeof(int), &pxl));
|
|
|
+ UFO_RESOURCES_CHECK_CLERR(clSetKernelArg( k_brep, 6, sizeof(int), &pyt));
|
|
|
+
|
|
|
+ UFO_RESOURCES_CHECK_CLERR(clEnqueueNDRangeKernel(cmd_queue, k_brep,
|
|
|
+ 2, NULL, priv->global_work_size_large, NULL,
|
|
|
+ 0, NULL, NULL));
|
|
|
+ }
|
|
|
+
|
|
|
+ return TRUE;
|
|
|
+}
|
|
|
+
|
|
|
+static void
|
|
|
+ufo_padding_2d_task_setup (UfoTask *task,
|
|
|
+ UfoResources *resources,
|
|
|
+ GError **error)
|
|
|
+{
|
|
|
+ UfoPadding2DTaskPrivate *priv;
|
|
|
+
|
|
|
+ priv = UFO_PADDING_2D_TASK_GET_PRIVATE (task);
|
|
|
+
|
|
|
+ priv->kernel_iconst = ufo_resources_get_kernel (resources, "padding_2d.cl", "padding_2d_init_const", error);
|
|
|
+ priv->kernel_cpyimg = ufo_resources_get_kernel (resources, "padding_2d.cl", "padding_2d_copy_in", error);
|
|
|
+ priv->kernel_brep = ufo_resources_get_kernel (resources, "padding_2d.cl", "padding_2d_brep", error);
|
|
|
+}
|
|
|
+
|
|
|
+static void
|
|
|
+ufo_padding_2d_task_get_requisition (UfoTask *task,
|
|
|
+ UfoBuffer **inputs,
|
|
|
+ UfoRequisition *requisition)
|
|
|
+{
|
|
|
+ UfoPadding2DTaskPrivate *priv;
|
|
|
+ UfoRequisition in_req;
|
|
|
+
|
|
|
+ priv = UFO_PADDING_2D_TASK_GET_PRIVATE (task);
|
|
|
+ ufo_buffer_get_requisition (inputs[0], &in_req);
|
|
|
+
|
|
|
+ priv->in_width = in_req.dims[0];
|
|
|
+ priv->in_height = in_req.dims[1];
|
|
|
+
|
|
|
+ requisition->n_dims = 2;
|
|
|
+ requisition->dims[0] = priv->out_width = priv->xl + priv->in_width + priv->xr;
|
|
|
+ requisition->dims[1] = priv->out_height = priv->yt + priv->in_height + priv->yb;
|
|
|
+
|
|
|
+ priv->global_work_size_small[0] = (size_t) priv->in_width;
|
|
|
+ priv->global_work_size_small[1] = (size_t) priv->in_height;
|
|
|
+ priv->global_work_size_large[0] = requisition->dims[0];
|
|
|
+ priv->global_work_size_large[1] = requisition->dims[1];
|
|
|
+}
|
|
|
+
|
|
|
+static void
|
|
|
+ufo_padding_2d_task_get_structure (UfoTask *task,
|
|
|
+ guint *n_inputs,
|
|
|
+ guint **n_dims,
|
|
|
+ UfoTaskMode *mode)
|
|
|
+{
|
|
|
+ UfoPadding2DTaskPrivate *priv;
|
|
|
+
|
|
|
+ priv = UFO_PADDING_2D_TASK_GET_PRIVATE (task);
|
|
|
+ *mode = UFO_TASK_MODE_SINGLE;
|
|
|
+ *n_inputs = 1;
|
|
|
+ *n_dims = g_new0 (guint, 1);
|
|
|
+ (*n_dims)[0] = 2;
|
|
|
+}
|
|
|
+
|
|
|
+static UfoNode *
|
|
|
+ufo_padding_2d_task_copy_real (UfoNode *node,
|
|
|
+ GError **error)
|
|
|
+{
|
|
|
+ UfoPadding2DTask *orig;
|
|
|
+ UfoPadding2DTask *copy;
|
|
|
+
|
|
|
+ orig = UFO_PADDING_2D_TASK (node);
|
|
|
+ copy = UFO_PADDING_2D_TASK (ufo_padding_2d_task_new ());
|
|
|
+
|
|
|
+ copy->priv->xl = orig->priv->xl;
|
|
|
+ copy->priv->xr = orig->priv->xr;
|
|
|
+ copy->priv->yb = orig->priv->yb;
|
|
|
+ copy->priv->yt = orig->priv->yt;
|
|
|
+ copy->priv->mode = orig->priv->mode;
|
|
|
+ copy->priv->pconst = orig->priv->pconst;
|
|
|
+
|
|
|
+ return UFO_NODE (copy);
|
|
|
+}
|
|
|
+
|
|
|
+static gboolean
|
|
|
+ufo_padding_2d_task_equal_real (UfoNode *n1,
|
|
|
+ UfoNode *n2)
|
|
|
+{
|
|
|
+ g_return_val_if_fail (UFO_IS_PADDING_2D_TASK (n1) && UFO_IS_PADDING_2D_TASK (n2), FALSE);
|
|
|
+ return TRUE;
|
|
|
+}
|
|
|
+
|
|
|
+static void
|
|
|
+ufo_padding_2d_task_finalize (GObject *object)
|
|
|
+{
|
|
|
+ UfoPadding2DTaskPrivate *priv;
|
|
|
+
|
|
|
+ priv = UFO_PADDING_2D_TASK_GET_PRIVATE (object);
|
|
|
+
|
|
|
+ G_OBJECT_CLASS (ufo_padding_2d_task_parent_class)->finalize (object);
|
|
|
+}
|
|
|
+
|
|
|
+static void
|
|
|
+ufo_task_interface_init (UfoTaskIface *iface)
|
|
|
+{
|
|
|
+ iface->setup = ufo_padding_2d_task_setup;
|
|
|
+ iface->get_requisition = ufo_padding_2d_task_get_requisition;
|
|
|
+ iface->get_structure = ufo_padding_2d_task_get_structure;
|
|
|
+}
|
|
|
+
|
|
|
+static void
|
|
|
+ufo_gpu_task_interface_init (UfoGpuTaskIface *iface)
|
|
|
+{
|
|
|
+ iface->process = ufo_padding_2d_task_process;
|
|
|
+}
|
|
|
+
|
|
|
+static void
|
|
|
+ufo_padding_2d_task_set_property (GObject *object,
|
|
|
+ guint property_id,
|
|
|
+ const GValue *value,
|
|
|
+ GParamSpec *pspec)
|
|
|
+{
|
|
|
+ UfoPadding2DTaskPrivate *priv = UFO_PADDING_2D_TASK_GET_PRIVATE (object);
|
|
|
+
|
|
|
+ switch (property_id) {
|
|
|
+ case PROP_XL:
|
|
|
+ priv->xl = g_value_get_uint(value);
|
|
|
+ break;
|
|
|
+ case PROP_XR:
|
|
|
+ priv->xr = g_value_get_uint(value);
|
|
|
+ break;
|
|
|
+ case PROP_YT:
|
|
|
+ priv->yt = g_value_get_uint(value);
|
|
|
+ break;
|
|
|
+ case PROP_YB:
|
|
|
+ priv->yb = g_value_get_uint(value);
|
|
|
+ break;
|
|
|
+ case PROP_MODE:
|
|
|
+ if (!g_strcmp0(g_value_get_string(value), "zero"))
|
|
|
+ priv->mode = PADDING_ZERO;
|
|
|
+ else if (!g_strcmp0(g_value_get_string(value), "const"))
|
|
|
+ priv->mode = PADDING_CONST;
|
|
|
+ else if (!g_strcmp0(g_value_get_string(value), "gavg"))
|
|
|
+ priv->mode = PADDING_GAVG;
|
|
|
+ else if (!g_strcmp0(g_value_get_string(value), "brep"))
|
|
|
+ priv->mode = PADDING_BREP;
|
|
|
+ else
|
|
|
+ G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec);
|
|
|
+ break;
|
|
|
+ case PROP_PCONST:
|
|
|
+ priv->pconst = (float) g_value_get_double(value);
|
|
|
+ break;
|
|
|
+ default:
|
|
|
+ G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
|
|
|
+ break;
|
|
|
+ }
|
|
|
+}
|
|
|
+
|
|
|
+static void
|
|
|
+ufo_padding_2d_task_get_property (GObject *object,
|
|
|
+ guint property_id,
|
|
|
+ GValue *value,
|
|
|
+ GParamSpec *pspec)
|
|
|
+{
|
|
|
+ UfoPadding2DTaskPrivate *priv = UFO_PADDING_2D_TASK_GET_PRIVATE (object);
|
|
|
+
|
|
|
+ switch (property_id) {
|
|
|
+ case PROP_XL:
|
|
|
+ g_value_set_uint(value, priv->xl);
|
|
|
+ break;
|
|
|
+ case PROP_XR:
|
|
|
+ g_value_set_uint(value, priv->xr);
|
|
|
+ break;
|
|
|
+ case PROP_YT:
|
|
|
+ g_value_set_uint(value, priv->yt);
|
|
|
+ break;
|
|
|
+ case PROP_YB:
|
|
|
+ g_value_set_uint(value, priv->yb);
|
|
|
+ break;
|
|
|
+ case PROP_MODE:
|
|
|
+ switch (priv->mode) {
|
|
|
+ case PADDING_ZERO:
|
|
|
+ g_value_set_string(value, "zero");
|
|
|
+ break;
|
|
|
+ case PADDING_CONST:
|
|
|
+ g_value_set_string(value, "const");
|
|
|
+ break;
|
|
|
+ case PADDING_GAVG:
|
|
|
+ g_value_set_string(value, "gavg");
|
|
|
+ break;
|
|
|
+ case PADDING_BREP:
|
|
|
+ g_value_set_string(value, "brep");
|
|
|
+ break;
|
|
|
+ default:
|
|
|
+ G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec);
|
|
|
+ break;
|
|
|
+ }
|
|
|
+ break;
|
|
|
+ case PROP_PCONST:
|
|
|
+ g_value_set_double(value, priv->pconst);
|
|
|
+ break;
|
|
|
+ default:
|
|
|
+ G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec);
|
|
|
+ break;
|
|
|
+ }
|
|
|
+}
|
|
|
+
|
|
|
+static void
|
|
|
+ufo_padding_2d_task_class_init (UfoPadding2DTaskClass *klass)
|
|
|
+{
|
|
|
+ GObjectClass *oclass;
|
|
|
+ UfoNodeClass *node_class;
|
|
|
+
|
|
|
+ oclass = G_OBJECT_CLASS (klass);
|
|
|
+ node_class = UFO_NODE_CLASS (klass);
|
|
|
+
|
|
|
+ oclass->finalize = ufo_padding_2d_task_finalize;
|
|
|
+ oclass->set_property = ufo_padding_2d_task_set_property;
|
|
|
+ oclass->get_property = ufo_padding_2d_task_get_property;
|
|
|
+
|
|
|
+ properties[PROP_XL] =
|
|
|
+ g_param_spec_uint("xl",
|
|
|
+ "Number of additional pixel on the left hand image side",
|
|
|
+ "Number of additional pixel on the left hand image side",
|
|
|
+ 0, 16384, 1,
|
|
|
+ G_PARAM_READWRITE);
|
|
|
+
|
|
|
+ properties[PROP_XR] =
|
|
|
+ g_param_spec_uint("xr",
|
|
|
+ "Number of additional pixel on the right hand image side",
|
|
|
+ "Number of additional pixel on the right hand image side",
|
|
|
+ 0, 16384, 1,
|
|
|
+ G_PARAM_READWRITE);
|
|
|
+
|
|
|
+ properties[PROP_YT] =
|
|
|
+ g_param_spec_uint("yt",
|
|
|
+ "Number of additional pixel on the top image side",
|
|
|
+ "Number of additional pixel on the top image side",
|
|
|
+ 0, 16384, 1,
|
|
|
+ G_PARAM_READWRITE);
|
|
|
+
|
|
|
+ properties[PROP_YB] =
|
|
|
+ g_param_spec_uint("yb",
|
|
|
+ "Number of additional pixel on the bottom image side",
|
|
|
+ "Number of additional pixel on the bottom image side",
|
|
|
+ 0, 16384, 1,
|
|
|
+ G_PARAM_READWRITE);
|
|
|
+
|
|
|
+ properties[PROP_MODE] =
|
|
|
+ g_param_spec_string("mode",
|
|
|
+ "Padding mode can be 'zero', 'const', 'gavg' or 'brep' ",
|
|
|
+ "Padding mode can be 'zero', 'const', 'gavg' or 'brep' ",
|
|
|
+ "zero",
|
|
|
+ G_PARAM_READWRITE);
|
|
|
+
|
|
|
+ properties[PROP_PCONST] =
|
|
|
+ g_param_spec_double("pconst",
|
|
|
+ "Padding constant",
|
|
|
+ "Padding constant",
|
|
|
+ -320000.0,
|
|
|
+ 320000.0,
|
|
|
+ 0.0,
|
|
|
+ G_PARAM_READWRITE);
|
|
|
+
|
|
|
+ for (guint i = PROP_0 + 1; i < N_PROPERTIES; i++)
|
|
|
+ g_object_class_install_property (oclass, i, properties[i]);
|
|
|
+
|
|
|
+ node_class->copy = ufo_padding_2d_task_copy_real;
|
|
|
+ node_class->equal = ufo_padding_2d_task_equal_real;
|
|
|
+
|
|
|
+ g_type_class_add_private(klass, sizeof(UfoPadding2DTaskPrivate));
|
|
|
+}
|
|
|
+
|
|
|
+static void
|
|
|
+ufo_padding_2d_task_init (UfoPadding2DTask *self)
|
|
|
+{
|
|
|
+ UfoPadding2DTaskPrivate *priv;
|
|
|
+ self->priv = priv = UFO_PADDING_2D_TASK_GET_PRIVATE (self);
|
|
|
+
|
|
|
+ priv->xl = 1;
|
|
|
+ priv->xr = 1;
|
|
|
+ priv->yt = 1;
|
|
|
+ priv->yb = 1;
|
|
|
+
|
|
|
+ priv->mode = PADDING_ZERO;
|
|
|
+ priv->pconst = 0.0;
|
|
|
+
|
|
|
+ priv->kernel_iconst = NULL;
|
|
|
+ priv->kernel_cpyimg = NULL;
|
|
|
+ priv->kernel_brep = NULL;
|
|
|
+}
|
Matthias Vogelgesang