Merge pull request #169 from ufo-kit/allow_more_input_nodes

Allow more input nodes

CMakeLists.txt

@@ -8,7 +8,7 @@ if (POLICY CMP0054)
 cmake_policy(SET CMP0054 OLD)
 endif ()
 
-project(ufo C)
+project(ufo)
 
 set(TARNAME "libufo")
 
@@ -47,6 +47,7 @@ endif()
 #{{{ Options
 option(WITH_TESTS "Build test suite" ON)
 option(WITH_DEPRECATED_OPENCL_1_1_API "Build with deprecated OpenCL 1.1 API" ON)
+set(UFO_MAX_INPUT_NODES "64" CACHE STRING "Maximum number of allowed input nodes for a task")
 
 if (WITH_DEPRECATED_OPENCL_1_1_API)
     add_definitions ("-DCL_USE_DEPRECATED_OPENCL_1_1_APIS")

config.h.in

@@ -1,5 +1,6 @@
 #cmakedefine WITH_PYTHON    1
 #cmakedefine HAVE_VIENNACL  1
+#cmakedefine UFO_MAX_INPUT_NODES ${UFO_MAX_INPUT_NODES}
 #define UFO_PLUGIN_DIR  "${CMAKE_INSTALL_PREFIX}/${CMAKE_INSTALL_PLUGINDIR}"
 #define UFO_KERNEL_DIR  "${CMAKE_INSTALL_PREFIX}/${CMAKE_INSTALL_KERNELDIR}"
 #define UFO_VERSION     "${UFO_VERSION}"

config.h.meson.in

@@ -6,3 +6,4 @@
 #mesondefine GLIB_VERSION_MIN_REQUIRED
 #mesondefine GLIB_VERSION_MAX_ALLOWED
 #mesondefine CL_TARGET_OPENCL_VERSION
+#mesondefine UFO_MAX_INPUT_NODES

meson.build

@@ -1,4 +1,4 @@
-project('ufo', 'c',
+project('ufo', ['c', 'cpp'],
     version: '0.16.0',
     default_options: ['c_std=gnu99'],
 )
@@ -65,6 +65,7 @@ conf.set_quoted('UFO_VERSION', version)
 conf.set('GLIB_VERSION_MIN_REQUIRED', 'GLIB_VERSION_2_38')
 conf.set('GLIB_VERSION_MAX_ALLOWED', 'GLIB_VERSION_2_38')
 conf.set('CL_TARGET_OPENCL_VERSION', '120')
+conf.set('UFO_MAX_INPUT_NODES', get_option('ufo_max_input_nodes'))
 
 if python_dep.found()
   conf.set('WITH_PYTHON', true)

meson_options.txt

@@ -9,3 +9,7 @@ option('introspection',
 option('bashcompletiondir',
     type: 'string',
     description: 'directory for bash complection scripts ["no" disables]')
+
+option('ufo_max_input_nodes',
+    type: 'integer', value: 64,
+    description: 'Maximum number of allowed input nodes for a task')

tests/unit/CMakeLists.txt

@@ -6,6 +6,7 @@ set(TEST_SRCS
     test-graph.c
     test-node.c
     test-profiler.c
+    test-max-input-nodes.cpp
     )
 
 set(SUITE_BIN "test-suite")

tests/unit/meson.build

@@ -4,6 +4,7 @@ sources = [
     'test-graph.c',
     'test-node.c',
     'test-profiler.c',
+    'test-max-input-nodes.cpp'
 ]
 
 test('unit tests',

tests/unit/test-max-input-nodes.cpp

@@ -0,0 +1,148 @@
+/*
+ * Copyright (C) 2011-2013 Karlsruhe Institute of Technology
+ *
+ * This file is part of Ufo.
+ *
+ * This library is free software: you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation, either
+ * version 3 of the License, or (at your option) any later version.
+ *
+ * This library is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with this library.  If not, see <http://www.gnu.org/licenses/>.
+ */
+
+#include <glib.h>
+#include "config.h"
+#include <ufo/ufo.h>
+
+#include <string>
+#include <vector>
+#include <stdexcept>
+#include <iostream>
+
+/**
+ * Build a opencl-kernel that sums *num_inputs* streams slice wise.
+ *
+ * @param num_inputs Number of input-streams
+ * @return std::string containing an opencl-kernel
+ */
+std::string build_kernel(unsigned int num_inputs){
+    if (num_inputs < 1 or num_inputs > 99999)
+        throw std::invalid_argument("num_inputs must be between 0 and 99999.");
+
+    std::string kernel = "kernel void test_input(";
+    for(unsigned int i = 0; i < num_inputs; ++i){
+        kernel.append("global float *a");
+        char number[8];
+        sprintf(number, "%06i", i);
+        kernel.append(number);
+        kernel.append(", ");
+    }
+    kernel.append("global float *result){\n");
+    kernel.append("\tsize_t idx = get_global_id(1) * get_global_size(0) + get_global_id(0);\n");
+    kernel.append("\tresult[idx] = ");
+    for(unsigned int i = 0; i < num_inputs; ++i){
+        kernel.append("a");
+        char number[8];
+        sprintf(number, "%06i", i);
+        kernel.append(number);
+        kernel.append("[idx]");
+        if (i < num_inputs - 1) {
+            kernel.append(" + ");
+        }
+        else{
+            kernel.append(";\n");
+        }
+    }
+    kernel.append("}\n");
+    return kernel;
+}
+/**
+ * Creates a graph with *n* dummy-data inputs.
+ *
+ * The inputs are connected to a opencl-node, that sums the inputs slice wise.
+ * Then the summed data is put in a null-sink.
+ *
+ * @param n Number of inputs
+ */
+static void test_n_inputs(unsigned int n){
+    UfoTaskGraph *graph;
+    UfoBaseScheduler *scheduler;
+    UfoPluginManager *manager;
+
+#if !(GLIB_CHECK_VERSION(2, 36, 0))
+    g_type_init ();
+#endif
+
+    graph = UFO_TASK_GRAPH(ufo_task_graph_new());
+    manager = ufo_plugin_manager_new();
+    scheduler = ufo_scheduler_new();
+
+    auto opencl_kernel = ufo_plugin_manager_get_task(manager, "opencl", nullptr);
+    if (opencl_kernel == nullptr)
+        throw std::runtime_error("Can not create task 'opencl'.");
+    g_object_set (G_OBJECT (opencl_kernel),
+                  "source", build_kernel(n).c_str(),
+                  "kernel", "test_input",
+                  NULL);
+
+    auto sink = ufo_plugin_manager_get_task(manager, "null", nullptr);
+    if (sink == nullptr)
+        throw std::runtime_error("Can not create task 'null'.");
+
+    std::vector<UfoTaskNode *> readers;
+    readers.resize(n);
+    std::size_t input_node_counter = 0;
+    for (auto& reader: readers) {
+        reader = ufo_plugin_manager_get_task(manager, "dummy-data", nullptr);
+        if (reader == nullptr)
+            throw std::runtime_error("Can not create task 'dummy-data'.");
+        g_object_set (G_OBJECT (reader),
+                      "width", 256,
+                      "height", 256,
+                      "number", 100,
+                      NULL);
+        ufo_task_graph_connect_nodes_full(graph, reader, opencl_kernel, input_node_counter);
+        input_node_counter++;
+    }
+
+    ufo_task_graph_connect_nodes(graph, opencl_kernel, sink);
+    ufo_base_scheduler_run(scheduler, graph, nullptr);
+
+    /* Destroy all objects */
+    for (const auto& reader: readers) {
+        g_object_unref(reader);
+    }
+    g_object_unref(graph);
+    g_object_unref(scheduler);
+    g_object_unref(manager);
+    g_object_unref(opencl_kernel);
+    g_object_unref(sink);
+
+}
+
+/**
+ * Tests a graph with a node, featuring 1 to UFO_MAX_INPUT_NODES.
+ */
+static void test_max_inputs() {
+    for(int i = 1; i <= UFO_MAX_INPUT_NODES; ++i){
+        try {
+            test_n_inputs(i);
+        }
+        catch (const std::exception& e){
+            std::cout << e.what() << std::endl;
+            g_assert(false);
+        }
+    }
+}
+
+
+extern "C" void test_add_max_input_nodes(void) {
+        g_test_add_func("/no-opencl/scheduler/max_input_nodes", test_max_inputs);
+}

tests/unit/test-suite.c

@@ -46,6 +46,7 @@ int main(int argc, char *argv[])
     test_add_graph ();
     test_add_profiler ();
     test_add_node ();
+    test_add_max_input_nodes();
 
     g_test_run();
 

tests/unit/test-suite.h

@@ -24,5 +24,6 @@ void test_add_buffer (void);
 void test_add_graph (void);
 void test_add_node (void);
 void test_add_profiler (void);
+void test_add_max_input_nodes(void);
 
 #endif

ufo/ufo-base-scheduler.c

@@ -71,6 +71,7 @@ enum {
     PROP_ENABLE_TRACING,
     PROP_TIMESTAMPS,
     PROP_TIME,
+    PROP_MAX_INPUT_NODES,
     N_PROPERTIES,
 };
 
@@ -299,6 +300,10 @@ ufo_base_scheduler_get_property (GObject *object,
             g_value_set_double (value, priv->time);
             break;
 
+        case PROP_MAX_INPUT_NODES:
+            g_value_set_uint(value, UFO_MAX_INPUT_NODES);
+            break;
+
         default:
             G_OBJECT_WARN_INVALID_PROPERTY_ID(object, property_id, pspec);
             break;
@@ -406,6 +411,13 @@ ufo_base_scheduler_class_init (UfoBaseSchedulerClass *klass)
                               0.0, G_MAXDOUBLE, 0.0,
                               G_PARAM_READABLE);
 
+    properties[PROP_MAX_INPUT_NODES] =
+        g_param_spec_uint("max_input_nodes",
+                          "Maximum inputs per task",
+                          "Maximum inputs per task",
+                          1, G_MAXUINT, UFO_MAX_INPUT_NODES,
+                          G_PARAM_READABLE);
+
     for (guint i = PROP_0 + 1; i < N_PROPERTIES; i++)
         g_object_class_install_property (oclass, i, properties[i]);
 

ufo/ufo-scheduler.c

@@ -307,15 +307,15 @@ check_target_connections (UfoTaskGraph *graph,
 {
     GList *predecessors;
     GList *it;
-    guint16 connection_bitmap;
-    guint16 mask;
+    gboolean connections[UFO_MAX_INPUT_NODES];
     gboolean result = TRUE;
 
     if (n_inputs == 0)
         return TRUE;
 
     predecessors = ufo_graph_get_predecessors (UFO_GRAPH (graph), target);
-    connection_bitmap = 0;
+
+    gboolean all_inputs_connected = TRUE;
 
     /* Check all edges and enable bit number for edge label */
     g_list_for (predecessors, it) {
@@ -325,14 +325,20 @@ check_target_connections (UfoTaskGraph *graph,
         label = ufo_graph_get_edge_label (UFO_GRAPH (graph),
                                           UFO_NODE (it->data), target);
         input = GPOINTER_TO_INT (label);
-        g_assert (input >= 0 && input < 16);
-        connection_bitmap |= 1 << input;
-    }
 
-    mask = (1 << n_inputs) - 1;
+        if (input > UFO_MAX_INPUT_NODES) {
+            g_set_error(error, UFO_SCHEDULER_ERROR, UFO_SCHEDULER_ERROR_SETUP,
+                        "Number of inputs exceeds %i (UFO_MAX_INPUT_NODES). This value can be defined in the build settings.",
+                        UFO_MAX_INPUT_NODES);
+        }
+
+        g_assert (input >= 0 && input < UFO_MAX_INPUT_NODES);
+        connections[input] = TRUE;
+        all_inputs_connected &= connections[input];
+    }
 
     /* Check if mask matches what we have */
-    if ((mask & connection_bitmap) != mask) {
+    if (!all_inputs_connected) {
         g_set_error (error, UFO_SCHEDULER_ERROR, UFO_SCHEDULER_ERROR_SETUP,
                      "Not all inputs of `%s' are connected",
                      ufo_task_node_get_plugin_name (UFO_TASK_NODE (target)));

ufo/ufo-task-node.c

@@ -49,9 +49,9 @@ struct _UfoTaskNodePrivate {
     UfoNode         *proc_node;
     UfoGroup        *out_group;
     UfoProfiler     *profiler;
-    GList           *in_groups[16];
-    GList           *current[16];
-    gint             n_expected[16];
+    GList           *in_groups[UFO_MAX_INPUT_NODES];
+    GList           *current[UFO_MAX_INPUT_NODES];
+    gint             n_expected[UFO_MAX_INPUT_NODES];
     guint            index;
     guint            total;
     guint            num_processed;
@@ -133,7 +133,7 @@ ufo_task_node_set_num_expected (UfoTaskNode *node,
                                 gint n_expected)
 {
     g_return_if_fail (UFO_IS_TASK_NODE (node));
-    g_return_if_fail (pos < 16);
+    g_return_if_fail (pos < UFO_MAX_INPUT_NODES);
     node->priv->n_expected[pos] = n_expected;
 }
 
@@ -142,7 +142,7 @@ ufo_task_node_get_num_expected (UfoTaskNode *node,
                                 guint pos)
 {
     g_return_val_if_fail (UFO_IS_TASK_NODE (node), 0);
-    g_return_val_if_fail (pos < 16, 0);
+    g_return_val_if_fail (pos < UFO_MAX_INPUT_NODES, 0);
     return node->priv->n_expected[pos];
 }
 
@@ -197,7 +197,7 @@ ufo_task_node_reset (UfoTaskNode *node)
     priv->out_group = NULL;
     priv->proc_node = NULL;
 
-    for (guint i = 0; i < 16; i++) {
+    for (guint i = 0; i < UFO_MAX_INPUT_NODES; i++) {
         g_list_free (priv->in_groups[i]);
         priv->in_groups[i] = NULL;
     }
@@ -219,7 +219,7 @@ ufo_task_node_get_current_in_group (UfoTaskNode *node,
                                     guint pos)
 {
     g_return_val_if_fail (UFO_IS_TASK_NODE (node), NULL);
-    g_assert (pos < 16);
+    g_assert (pos < UFO_MAX_INPUT_NODES);
     g_assert (node->priv->current[pos] != NULL);
     return UFO_GROUP (node->priv->current[pos]->data);
 }
@@ -329,7 +329,7 @@ ufo_task_node_copy (UfoNode *node,
 
     copy->priv->pattern = orig->priv->pattern;
 
-    for (guint i = 0; i < 16; i++)
+    for (guint i = 0; i < UFO_MAX_INPUT_NODES; i++)
         copy->priv->n_expected[i] = orig->priv->n_expected[i];
 
     ufo_task_node_set_plugin_name (copy, orig->priv->plugin);
@@ -433,7 +433,7 @@ ufo_task_node_init (UfoTaskNode *self)
     self->priv->num_processed = 0;
     self->priv->profiler = ufo_profiler_new ();
 
-    for (guint i = 0; i < 16; i++) {
+    for (guint i = 0; i < UFO_MAX_INPUT_NODES; i++) {
         self->priv->in_groups[i] = NULL;
         self->priv->current[i] = NULL;
         self->priv->n_expected[i] = -1;