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

#include <gmodule.h>
#ifdef __APPLE__
#include <OpenCL/cl.h>
#else
#include <CL/cl.h>
#endif

#include <ufo/ufo-resource-manager.h>
#include <ufo/ufo-filter.h>
#include <ufo/ufo-buffer.h>
#include "ufo-filter-padding-2d.h"

/**
 * SECTION:ufo-filter-padding-2d
 * @Short_description:
 * @Title: padding2d
 *
 * Detailed description.
 */

typedef enum {
    PADDING_ZERO = 0,
    PADDING_CONST,
    PADDING_GAVG,
    PADDING_BREP
} PaddingMode;


struct _UfoFilterPadding2DPrivate
{
    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;
};

G_DEFINE_TYPE(UfoFilterPadding2D, ufo_filter_padding_2d, UFO_TYPE_FILTER)

#define UFO_FILTER_PADDING_2D_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_FILTER_PADDING_2D, UfoFilterPadding2DPrivate))

enum {
    PROP_0 = 0,
    PROP_XL,
    PROP_XR,
    PROP_YT,
    PROP_YB,
    PROP_MODE,
    PROP_PCONST,
    N_PROPERTIES
};

static GParamSpec *padding_2d_properties[N_PROPERTIES] = { NULL, };


static void
ufo_filter_padding_2d_initialize (UfoFilter *filter, UfoBuffer *params[], guint **dim_sizes, GError **error)
{
    UfoFilterPadding2DPrivate *priv = UFO_FILTER_PADDING_2D_GET_PRIVATE(filter);
    UfoResourceManager *manager = ufo_filter_get_resource_manager(filter);
    GError *tmp_error = NULL;

    priv->kernel_iconst = ufo_resource_manager_get_kernel(manager, "padding_2d.cl", "padding_2d_init_const", &tmp_error);
    priv->kernel_cpyimg = ufo_resource_manager_get_kernel(manager, "padding_2d.cl", "padding_2d_copy_in", &tmp_error);
    priv->kernel_brep   = ufo_resource_manager_get_kernel(manager, "padding_2d.cl", "padding_2d_brep", &tmp_error);

    if (tmp_error != NULL) {
        g_propagate_error (error, tmp_error);
        return;
    }

    ufo_buffer_get_2d_dimensions (params[0], &priv->in_width, &priv->in_height);
    dim_sizes[0][0] = priv->out_width  = priv->xl + priv->in_width + priv->xr;
    dim_sizes[0][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] = (size_t) dim_sizes[0][0];
    priv->global_work_size_large[1] = (size_t) dim_sizes[0][1];
}

/*
 * This is the main method in which the filter processes one buffer after
 * another.
 */
static void
ufo_filter_padding_2d_process_gpu(UfoFilter *filter, UfoBuffer *input[], UfoBuffer *output[], GError **error)
{
    g_return_if_fail (UFO_IS_FILTER (filter));
    cl_command_queue command_queue = ufo_filter_get_command_queue (filter);
    UfoFilterPadding2DPrivate *priv = UFO_FILTER_PADDING_2D_GET_PRIVATE(filter);

    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;

    // init padding value  for different modes
    if (mode == PADDING_ZERO)
        pval = 0.0;

    cl_event event;
    cl_mem input_mem  = (cl_mem) ufo_buffer_get_device_array (input[0], command_queue);
    cl_mem output_mem = (cl_mem) ufo_buffer_get_device_array (output[0], command_queue);

    if (mode == PADDING_GAVG) {
        float *indata = ufo_buffer_get_host_array(input[0], command_queue);
        float sum = 0;
        guint psz = ixs * iys;

        for (guint i =0; i < psz; i++) 
            sum += indata[i];

        pval = sum / (float)psz;
    }

    // processing itself
    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
        CHECK_OPENCL_ERROR(clSetKernelArg( k_iconst, 0, sizeof(cl_mem), (void *) &output_mem));
        CHECK_OPENCL_ERROR(clSetKernelArg( k_iconst, 1, sizeof(int),   &oxs));
        CHECK_OPENCL_ERROR(clSetKernelArg( k_iconst, 2, sizeof(float), &pval));

        CHECK_OPENCL_ERROR(clEnqueueNDRangeKernel(command_queue,  k_iconst,
                    2, NULL, priv->global_work_size_large, NULL,
                    0, NULL, &event));
        //ufo_buffer_attach_event(output, event);

        /// copy old image
        CHECK_OPENCL_ERROR(clSetKernelArg(  k_cpyimg, 0, sizeof(cl_mem), (void *) &input_mem));
        CHECK_OPENCL_ERROR(clSetKernelArg(  k_cpyimg, 1, sizeof(cl_mem), (void *) &output_mem));
        CHECK_OPENCL_ERROR(clSetKernelArg(  k_cpyimg, 2, sizeof(int),   &ixs));
        CHECK_OPENCL_ERROR(clSetKernelArg(  k_cpyimg, 3, sizeof(int),   &oxs));
        CHECK_OPENCL_ERROR(clSetKernelArg(  k_cpyimg, 4, sizeof(int),   &pxl));
        CHECK_OPENCL_ERROR(clSetKernelArg(  k_cpyimg, 5, sizeof(int),   &pyt));

        //cl_event event2;
        CHECK_OPENCL_ERROR(clEnqueueNDRangeKernel(command_queue,  k_cpyimg,
                    2, NULL, priv->global_work_size_small, NULL,
                    0, NULL, &event));
        //ufo_buffer_attach_event(output, event2);
    }

    if (mode == PADDING_BREP) {
        cl_kernel k_brep   = priv->kernel_brep;

        CHECK_OPENCL_ERROR(clSetKernelArg( k_brep, 0, sizeof(cl_mem), (void *) &input_mem));
        CHECK_OPENCL_ERROR(clSetKernelArg( k_brep, 1, sizeof(cl_mem), (void *) &output_mem));
        CHECK_OPENCL_ERROR(clSetKernelArg( k_brep, 2, sizeof(int),   &ixs));
        CHECK_OPENCL_ERROR(clSetKernelArg( k_brep, 3, sizeof(int),   &iys));
        CHECK_OPENCL_ERROR(clSetKernelArg( k_brep, 4, sizeof(int),   &oxs));
        CHECK_OPENCL_ERROR(clSetKernelArg( k_brep, 5, sizeof(int),   &pxl));
        CHECK_OPENCL_ERROR(clSetKernelArg( k_brep, 6, sizeof(int),   &pyt));

        CHECK_OPENCL_ERROR(clEnqueueNDRangeKernel(command_queue,  k_brep,
                    2, NULL, priv->global_work_size_large, NULL,
                    0, NULL, &event));
        // ufo_buffer_attach_event(output, event);
    }

    clFinish(command_queue); // synchro?
}

static void
ufo_filter_padding_2d_set_property (GObject *object, guint property_id, const GValue *value, GParamSpec *pspec)
{
    UfoFilterPadding2D *self = UFO_FILTER_PADDING_2D(object);

    switch (property_id) {
        case PROP_XL:
            self->priv->xl = g_value_get_uint(value);
            break;
        case PROP_XR:
            self->priv->xr = g_value_get_uint(value);
            break;
        case PROP_YT:
            self->priv->yt = g_value_get_uint(value);
            break;
        case PROP_YB:
            self->priv->yb = g_value_get_uint(value);
            break;
        case PROP_MODE:
            if (!g_strcmp0(g_value_get_string(value), "zero"))
                self->priv->mode = PADDING_ZERO;
            else if (!g_strcmp0(g_value_get_string(value), "const"))
                self->priv->mode = PADDING_CONST;
            else if (!g_strcmp0(g_value_get_string(value), "gavg"))
                self->priv->mode = PADDING_GAVG;
            else if (!g_strcmp0(g_value_get_string(value), "brep"))
                self->priv->mode = PADDING_BREP;
            else
                G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec);
            break;
        case PROP_PCONST:
            self->priv->pconst = (float) g_value_get_double(value);
            break;
        default:
            G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec);
            break;
    }
}

static void
ufo_filter_padding_2d_get_property (GObject *object, guint property_id, GValue *value, GParamSpec *pspec)
{
    UfoFilterPadding2D *self = UFO_FILTER_PADDING_2D(object);

    switch (property_id) {
        case PROP_XL:
            g_value_set_uint(value, self->priv->xl);
            break;
        case PROP_XR:
            g_value_set_uint(value, self->priv->xr);
            break;
        case PROP_YT:
            g_value_set_uint(value, self->priv->yt);
            break;
        case PROP_YB:
            g_value_set_uint(value, self->priv->yb);
            break;
        case PROP_MODE:
            switch (self->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, self->priv->pconst);
            break;
        default:
            G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec);
            break;
    }
}

static void ufo_filter_padding_2d_class_init(UfoFilterPadding2DClass *klass)
{
    GObjectClass *gobject_class = G_OBJECT_CLASS(klass);
    UfoFilterClass *filter_class = UFO_FILTER_CLASS(klass);

    gobject_class->set_property = ufo_filter_padding_2d_set_property;
    gobject_class->get_property = ufo_filter_padding_2d_get_property;
    filter_class->initialize = ufo_filter_padding_2d_initialize;
    filter_class->process_gpu = ufo_filter_padding_2d_process_gpu;

    padding_2d_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);

    padding_2d_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);

    padding_2d_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);

    padding_2d_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);

    padding_2d_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);

    padding_2d_properties[PROP_PCONST] =
        g_param_spec_double("pconst",
                "Padding constant",
                "Padding constant",
                -320000.0,
                320000.0,
                0.0,
                G_PARAM_READWRITE);

    g_object_class_install_property(gobject_class, PROP_XL, padding_2d_properties[PROP_XL]);
    g_object_class_install_property(gobject_class, PROP_XR, padding_2d_properties[PROP_XR]);
    g_object_class_install_property(gobject_class, PROP_YT, padding_2d_properties[PROP_YT]);
    g_object_class_install_property(gobject_class, PROP_YB, padding_2d_properties[PROP_YB]);
    g_object_class_install_property(gobject_class, PROP_MODE, padding_2d_properties[PROP_MODE]);
    g_object_class_install_property(gobject_class, PROP_PCONST, padding_2d_properties[PROP_PCONST]);

    g_type_class_add_private(gobject_class, sizeof(UfoFilterPadding2DPrivate));
}

static void
ufo_filter_padding_2d_init(UfoFilterPadding2D *self)
{
    UfoFilterPadding2DPrivate *priv = self->priv = UFO_FILTER_PADDING_2D_GET_PRIVATE(self);
    UfoInputParameter input_params[] = {{2, UFO_FILTER_INFINITE_INPUT}};
    UfoOutputParameter output_params[] = {{2}};

    // set default values
    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;

    // register filter input and output
    ufo_filter_register_inputs (UFO_FILTER (self), 1, input_params);
    ufo_filter_register_outputs (UFO_FILTER (self), 1, output_params);
}

G_MODULE_EXPORT
UfoFilter *ufo_filter_plugin_new(void)
{
    return g_object_new(UFO_TYPE_FILTER_PADDING_2D, NULL);
}