Developing new task filters

UFO filters are simple shared objects that expose their GType and implement the UfoFilter class.

Writing a task in C

Writing a new UFO filter consists of filling out a pre-defined class structure. To avoid writing the GObject boild plate code, you can call

ufo-mkfilter AwesomeFoo

to generate the header and source file templates. The name must be a camel-cased version of your new filter.

You are now left with two files ufo-awesome-foo-task.c and ufo-awesome-foo-task.h. If you intend to distribute that filter with the main UFO filter distribution, copy these files to ufo-filters/src. If you are not depending on any third-party library you can just add the following line to the CMakeLists.txt file:

set(ufofilter_SRCS
    ...
    ufo-awesome-foo-task.c
    ...)

You can compile this as usual by typing

make

in the CMake build directory of ufo-filters.

Initializing filters

Regardless of filter type, the ufo_awesome_foo_task_init() method is the constructor of the filter object and the best place to setup all data that does not depend on any input data.

The ufo__awesome_foo_task_class_init() method on the other hand ist the class constructor that is used to override virtual methods by setting function pointers in each classes’ vtable.

You must override the following methods: ufo_awesome_task_get_num_inputs, ufo_awesome_task_get_num_dimensions, ufo_awesome_task_get_mode, ufo_awesome_task_setup, ufo_awesome_task_get_requisition, ufo_awesome_task_process and/or ufo_awesome_task_generate.

get_num_inputs, get_num_dimensions and get_mode are called by the run-time in order to determine how many inputs your task expects, which dimensions are allowed on each input and what processing mode your task runs

static guint
ufo_awesome_task_get_num_inputs (UfoTask *task)
{
    return 1;   /* We expect only one input */
}

static guint
ufo_awesome_task_get_num_dimensions (UfoTask *task, guint input)
{
    return 2;   /* We ignore "input" and always expect 2 dimensions */
}

static UfoTaskMode
ufo_awesome_task_get_mode (UfoTask *task)
{
    /* We process one item after another */
    return UFO_TASK_MODE_PROCESSOR;
}

The mode decides which functions of a task are called. Each task can provide a process function that takes input data and optionally writes output data and a generate function that does not take input data but writes data. Both functions return a boolean value to signal if data was produced or not (e.g. end of stream):

  • UFO_TASK_MODE_PROCESSOR: The task reads data and optionally writes data. For that it must implement process.
  • UFO_TASK_MODE_GENERATOR: The task only produces data (e.g. file readers) and must implement generate.
  • UFO_TASK_MODE_REDUCTOR: The tasks reads the input stream and produces another output stream. Reading is accomplished by implementing process whereas production is done by generate.

setup can be used to initialize data that depends on run-time resources like OpenCL contexts etc. This method is called only once

static void
ufo_awesome_task_setup (UfoTask *task,
                        UfoResources *resources,
                        GError **error)
{
    cl_context context;

    context = ufo_resources_get_context (resources);

    /*
       Do something with the context like allocating buffers or create
       kernels.
     */
}

On the other hand, get_requisition is called on each iteration right before process. It is used to determine which size an output buffer must have depending on the inputs. For this you must fill in the requisition structure correctly. If our output buffer needs to be as big as our input buffer we would specify

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

Finally, you have to override the process method

static gboolean
ufo_awesome_task_process (UfoTask *task,
                          UfoBuffer **inputs,
                          UfoBuffer *output,
                          UfoRequisition *requisition)
{
    UfoGpuNode *node;
    cl_command_queue cmd_queue;
    cl_mem host_in;
    cl_mem host_out;

    /* We have to know to which GPU device we are assigned to */
    node = UFO_GPU_NODE (ufo_task_node_get_proc_node (UFO_TASK_NODE (task)));

    /* Now, we can get the command queue */
    cmd_queue = ufo_gpu_node_get_cmd_queue (node);

    /* ... and get hold of the data */
    host_in = ufo_buffer_get_device_array (inputs[0], cmd_queue);
    host_out = ufo_buffer_get_device_array (output, cmd_queue);

    /* Call a kernel or do other meaningful work. */
}

Tasks can and will be copied to speed up the computation on multi-GPU systems. Any parameters that are accessible from the outside via a property are automatically copied by the run-time system. To copy private data that is only visible at the file scope, you have to override the UFO_NODE_CLASS method copy and copy the data yourself. This method is always called before setup so you can be assured to re-create your private data on the copied task.

Note

It is strongly encouraged that you export all your parameters as properties and re-build any internal data structures off of these parameters.

Additional source files

For modularity reasons, you might want to split your filter sources into different compilation units. In order to compile and link them against the correct library, add the following statements to the src/CMakeLists.txt file

set(awesome_foo_misc_SRCS foo.c bar.c baz.c)

in case your filter is still called AwesomeFoo. Notice, that the variable name matches the plugin name with underscores between the lower-cased letters.

Writing point-based OpenCL filters

For point-based image operations it is much faster to use the cl-plugin that writing a full-fledged C filter. We create a new file simple.cl, that contains a simple kernel that inverts our normalized input (you can silently ignore the scratch parameter for now):

kernel void invert(global float *input, global float *output)
{
    /* where are we? */
    int index = get_global_id(1) * get_global_size(0) + get_global_id(0);
    output[index] = 1.0f - input[index];
}

We wire this small kernel into this short Python script:

from gi.repository import Ufo

pm = Ufo.PluginManager()
reader = pm.get_filter('reader')
writer = pm.get_filter('writer')

# this filter applies the kernel
cl = pm.get_filter('opencl')
cl.set_properties(filename='simple.cl', kernel='invert')

g = Ufo.TaskGraph()
g.connect_nodes(reader, cl)
g.connect_nodes(cl, writer)

s = Ufo.Scheduler()
s.run(g)

For more information on how to write OpenCL kernels, consult the official OpenCL reference pages.

Reporting errors at run-time

From within a filter (or any library for that matter) never call functions such as exit() or abort(). This prevents the calling application from identification of the problem as well as recovery. Instead use the builtin GError infrastructure that – as a bonus – map nicely to exceptions in Python:

static void
ufo_awesome_task_setup (UfoTask *task,
                        UfoResources *resources,
                        GError **error)
{
    if (error_condition) {
        g_set_error (error, UFO_TASK_ERROR, UFO_TASK_ERROR_SETUP,
                     "Error because of condition");
        return;
    }
}

Note that g_set_error receives printf-style format strings which means you should be as specific as possible with the given error message.

The GObject property system

Wait until a property satisfies a condition

For some filters it could be important to not only wait until input buffers arrive but also properties change their values. For example, the back-projection should only start as soon as it is assigned a correct center-of-rotation. To implement this, we have to define a condition function that checks if a GValue representing the current property satisfies a certain condition

static gboolean is_larger_than_zero(GValue *value, gpointer user_data)
{
    return g_value_get_float(value) > 0.0f;
}

As the filter installed the properties it also knows which type it is and which g_value_get_*() function to call. Now, we wait until this conditions holds using ufo_filter_wait_until()

/* Somewhere in ufo_filter_process() */
ufo_filter_wait_until(self, properties[PROP_CENTER_OF_ROTATION],
        &is_larger_than_zero, NULL);

Warning

ufo_filter_wait_until() might block indefinitely when the condition function never returns TRUE.