/**
 * SECTION:ufo-filter-task
 * @Short_description: Process arbitrary Filter kernels
 * @Title: filter
 *
 * This module is used to load an arbitrary #UfoScaleTask:kernel from
 * #UfoScaleTask: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. #UfoScaleTask:num-dims must be changed, if the kernel
 * accesses either one or three dimensional index spaces.
 */

#ifdef __APPLE__
#include <Filter/cl.h>
#else
#include <CL/cl.h>
#endif
#include "ufo-scale-task.h"

struct _UfoScaleTaskPrivate {
    cl_kernel kernel;
    gfloat scale;
};

static void ufo_task_interface_init (UfoTaskIface *iface);

G_DEFINE_TYPE_WITH_CODE (UfoScaleTask, ufo_scale_task, UFO_TYPE_TASK_NODE,
                         G_IMPLEMENT_INTERFACE (UFO_TYPE_TASK,
                                                ufo_task_interface_init))

#define UFO_SCALE_TASK_GET_PRIVATE(obj) (G_TYPE_INSTANCE_GET_PRIVATE((obj), UFO_TYPE_SCALE_TASK, UfoScaleTaskPrivate))

enum {
    PROP_0,
    PROP_SCALE,
    N_PROPERTIES
};

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

UfoNode *
ufo_scale_task_new (void)
{
    return UFO_NODE (g_object_new (UFO_TYPE_SCALE_TASK, NULL));
}

static gboolean
ufo_scale_task_process (UfoTask *task,
                        UfoBuffer **inputs,
                        UfoBuffer *output,
                        UfoRequisition *requisition)
{
    UfoScaleTaskPrivate *priv;
    UfoGpuNode *node;
    cl_command_queue cmd_queue;
    cl_mem in_mem;
    cl_mem out_mem;

    priv = UFO_SCALE_TASK (task)->priv;
    node = UFO_GPU_NODE (ufo_task_node_get_proc_node (UFO_TASK_NODE (task)));
    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);

    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 0, sizeof (cl_mem), &in_mem));
    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 1, sizeof (cl_mem), &out_mem));
    UFO_RESOURCES_CHECK_CLERR (clSetKernelArg (priv->kernel, 2, sizeof (cl_float), &priv->scale));

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

    return TRUE;
}

static guint
ufo_scale_task_get_num_inputs (UfoTask *task)
{
    return 1;
}

static guint
ufo_scale_task_get_num_dimensions (UfoTask *task,
                                           guint input)
{
    g_return_val_if_fail (input == 0, 0);
    return 2;
}

static UfoTaskMode
ufo_scale_task_get_mode (UfoTask *task)
{
    return UFO_TASK_MODE_PROCESSOR | UFO_TASK_MODE_GPU;
}

static void
ufo_scale_task_setup (UfoTask *task,
                      UfoResources *resources,
                      GError **error)
{
    UfoScaleTaskPrivate *priv;

    priv = UFO_SCALE_TASK_GET_PRIVATE (task);

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

    if (priv->kernel != NULL)
        UFO_RESOURCES_CHECK_CLERR (clRetainKernel (priv->kernel));

}

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

static UfoNode *
ufo_scale_task_copy_real (UfoNode *node,
                          GError **error)
{
    UfoScaleTask *orig;
    UfoScaleTask *copy;

    orig = UFO_SCALE_TASK (node);
    copy = UFO_SCALE_TASK (ufo_scale_task_new ());

    g_object_set (G_OBJECT (copy),
                  "scale", orig->priv->scale,
                  NULL);

    return UFO_NODE (copy);
}

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

static void
ufo_scale_task_finalize (GObject *object)
{
    UfoScaleTaskPrivate *priv;

    priv = UFO_SCALE_TASK_GET_PRIVATE (object);

    if (priv->kernel) {
        clReleaseKernel (priv->kernel);
        priv->kernel = NULL;
    }

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

static void
ufo_task_interface_init (UfoTaskIface *iface)
{
    iface->setup = ufo_scale_task_setup;
    iface->get_requisition = ufo_scale_task_get_requisition;
    iface->get_num_inputs = ufo_scale_task_get_num_inputs;
    iface->get_num_dimensions = ufo_scale_task_get_num_dimensions;
    iface->get_mode = ufo_scale_task_get_mode;
    iface->process = ufo_scale_task_process;
}

static void
ufo_scale_task_set_property (GObject *object,
                             guint property_id,
                             const GValue *value,
                             GParamSpec *pspec)
{
    UfoScaleTaskPrivate *priv = UFO_SCALE_TASK_GET_PRIVATE (object);

    switch (property_id) {
        case PROP_SCALE:
            priv->scale = g_value_get_float (value);
            break;
        default:
            G_OBJECT_WARN_INVALID_PROPERTY_ID (object, property_id, pspec);
            break;
    }
}

static void
ufo_scale_task_get_property (GObject *object,
                              guint property_id,
                              GValue *value,
                              GParamSpec *pspec)
{
    UfoScaleTaskPrivate *priv = UFO_SCALE_TASK_GET_PRIVATE (object);

    switch (property_id) {
        case PROP_SCALE:
            g_value_set_float (value, priv->scale);
            break;
        default:
            G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec);
            break;
    }
}

static void
ufo_scale_task_class_init (UfoScaleTaskClass *klass)
{
    GObjectClass *oclass;
    UfoNodeClass *node_class;
    
    oclass = G_OBJECT_CLASS (klass);
    node_class = UFO_NODE_CLASS (klass);

    oclass->finalize = ufo_scale_task_finalize;
    oclass->set_property = ufo_scale_task_set_property;
    oclass->get_property = ufo_scale_task_get_property;

    properties[PROP_SCALE] =
        g_param_spec_float ("scale",
                            "Scale",
                            "Scale for each pixel",
                            -5.0, 10.0, 1.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_scale_task_copy_real;
    node_class->equal = ufo_scale_task_equal_real;

    g_type_class_add_private(klass, sizeof(UfoScaleTaskPrivate));
}

static void
ufo_scale_task_init (UfoScaleTask *self)
{
    UfoScaleTaskPrivate *priv;
    self->priv = priv = UFO_SCALE_TASK_GET_PRIVATE (self);
    priv->kernel = NULL;
}