lamino-filters/src/ufo-padding-2d-task.c

/**
 * 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;
}