lamino-filters/src/ufo-lamino-conv-task.c

/**
 * SECTION:ufo-filter-task
 * @Short_description: Process arbitrary Filter kernels
 * @Title: filter
 *
 * This module is used to load an arbitrary #UfoLaminoConvTask:kernel from
 * #UfoLaminoConvTask: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. #UfoLaminoConvTask: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-lamino-conv-task.h"

struct _UfoLaminoConvTaskPrivate {
    cl_kernel kernel;
    guint img_width;
    guint img_height;
};

static void ufo_task_interface_init (UfoTaskIface *iface);
static void ufo_gpu_task_interface_init (UfoGpuTaskIface *iface);

G_DEFINE_TYPE_WITH_CODE (UfoLaminoConvTask, ufo_lamino_conv_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_LAMINO_CONV_TASK_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_LAMINO_CONV_TASK, UfoLaminoConvTaskPrivate))

UfoNode *
ufo_lamino_conv_task_new (void)
{
    return UFO_NODE (g_object_new (UFO_TYPE_LAMINO_CONV_TASK, NULL));
}

static gboolean
ufo_lamino_conv_task_process (UfoGpuTask *task,
                              UfoBuffer **inputs,
                              UfoBuffer *output,
                              UfoRequisition *requisition,
                              UfoGpuNode *node)
{
    UfoLaminoConvTaskPrivate *priv;
    cl_command_queue cmd_queue;
    cl_mem in1_mem;
    cl_mem in2_mem;
    cl_mem out_mem;

    priv = UFO_LAMINO_CONV_TASK (task)->priv;

    cmd_queue = ufo_gpu_node_get_cmd_queue (node);
    in1_mem = ufo_buffer_get_device_array (inputs[0], cmd_queue);
    in2_mem = ufo_buffer_get_device_array (inputs[1], cmd_queue);
    out_mem = ufo_buffer_get_device_array (output, cmd_queue);

    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 0, sizeof(cl_mem), (void *) &in1_mem));
    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 1, sizeof(cl_mem), (void *) &in2_mem));
    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 2, sizeof(cl_mem), (void *) &out_mem));

    UFO_RESOURCES_CHECK_CLERR (clEnqueueNDRangeKernel (cmd_queue, priv->kernel,
                                                       2, NULL, requisition->dims, NULL,
                                                       0, NULL, NULL));

    return TRUE;
}

static void
ufo_lamino_conv_task_setup (UfoTask *task,
                           UfoResources *resources,
                           GError **error)
{
    UfoLaminoConvTaskPrivate *priv;

    priv = UFO_LAMINO_CONV_TASK_GET_PRIVATE (task);

    priv->kernel = ufo_resources_get_kernel (resources, "lamino_ft_conv.cl", "lamino_c", error);

    if (priv->kernel)
        clRetainKernel (priv->kernel);
}

static void
ufo_lamino_conv_task_get_requisition (UfoTask *task,
                                      UfoBuffer **inputs,
                                      UfoRequisition *requisition)
{
    ufo_buffer_get_requisition (inputs[0], requisition);
}

static void
ufo_lamino_conv_task_get_structure (UfoTask *task,
                                    guint *n_inputs,
                                    UfoInputParam **in_params,
                                    UfoTaskMode *mode)
{
    UfoLaminoConvTaskPrivate *priv;

    priv = UFO_LAMINO_CONV_TASK_GET_PRIVATE (task);
    *mode = UFO_TASK_MODE_SINGLE;
    *n_inputs = 2;
    *in_params = g_new0 (UfoInputParam, 2);
    (*in_params)[0].n_dims = 2;
    (*in_params)[0].n_expected = -1;
    (*in_params)[1].n_dims = 2;
    (*in_params)[1].n_expected = 1;
}

static UfoNode *
ufo_lamino_conv_task_copy_real (UfoNode *node,
                                GError **error)
{
    UfoLaminoConvTask *copy;
    copy = UFO_LAMINO_CONV_TASK (ufo_lamino_conv_task_new ());
    return UFO_NODE (copy);
}

static gboolean
ufo_lamino_conv_task_equal_real (UfoNode *n1,
                            UfoNode *n2)
{
    g_return_val_if_fail (UFO_IS_LAMINO_CONV_TASK (n1) && UFO_IS_LAMINO_CONV_TASK (n2), FALSE);
    return TRUE;
}

static void
ufo_lamino_conv_task_finalize (GObject *object)
{
    UfoLaminoConvTaskPrivate *priv;

    priv = UFO_LAMINO_CONV_TASK_GET_PRIVATE (object);

    G_OBJECT_CLASS (ufo_lamino_conv_task_parent_class)->finalize (object);
}

static void
ufo_task_interface_init (UfoTaskIface *iface)
{
    iface->setup = ufo_lamino_conv_task_setup;
    iface->get_requisition = ufo_lamino_conv_task_get_requisition;
    iface->get_structure = ufo_lamino_conv_task_get_structure;
}

static void
ufo_gpu_task_interface_init (UfoGpuTaskIface *iface)
{
    iface->process = ufo_lamino_conv_task_process;
}

static void
ufo_lamino_conv_task_set_property (GObject *object,
                                  guint property_id,
                                  const GValue *value,
                                  GParamSpec *pspec)
{
    switch (property_id) {
        default:
            G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
            break;
    }
}

static void
ufo_lamino_conv_task_get_property (GObject *object,
                              guint property_id,
                              GValue *value,
                              GParamSpec *pspec)
{
    switch (property_id) {
        default:
            G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec);
            break;
    }
}

static void
ufo_lamino_conv_task_class_init (UfoLaminoConvTaskClass *klass)
{
    GObjectClass *oclass;
    UfoNodeClass *node_class;
    
    oclass = G_OBJECT_CLASS (klass);
    node_class = UFO_NODE_CLASS (klass);

    oclass->finalize = ufo_lamino_conv_task_finalize;
    oclass->set_property = ufo_lamino_conv_task_set_property;
    oclass->get_property = ufo_lamino_conv_task_get_property;

    node_class->copy = ufo_lamino_conv_task_copy_real;
    node_class->equal = ufo_lamino_conv_task_equal_real;

    g_type_class_add_private(klass, sizeof(UfoLaminoConvTaskPrivate));
}

static void
ufo_lamino_conv_task_init (UfoLaminoConvTask *self)
{
    UfoLaminoConvTaskPrivate *priv;
    self->priv = priv = UFO_LAMINO_CONV_TASK_GET_PRIVATE (self);
}