Using multiple implementations of a codelet
===========================================

One may want to write multiple implementations of a codelet for a single type of device and let StarPU choose which one to run. As an example, we will show how to use SSE to scale a vector. The codelet can be written as follows:

     #include <xmmintrin.h>

     void scal_sse_func(void *buffers[], void *cl_arg)
     {
         float *vector = (float *) STARPU_VECTOR_GET_PTR(buffers[0]);
         unsigned int n = STARPU_VECTOR_GET_NX(buffers[0]);
         unsigned int n_iterations = n/4;
         if (n % 4 != 0)
             n_iterations++;

         __m128 *VECTOR = (__m128*) vector;
         __m128 factor __attribute__((aligned(16)));
         factor = _mm_set1_ps(*(float *) cl_arg);

         unsigned int i;
         for (i = 0; i < n_iterations; i++)
             VECTOR[i] = _mm_mul_ps(factor, VECTOR[i]);
     }

     struct starpu_codelet cl = {
         .where = STARPU_CPU,
         .cpu_funcs = { scal_cpu_func, scal_sse_func, NULL },
         .nbuffers = 1,
         .modes = { STARPU_RW }
     };

Schedulers which are multi-implementation aware (only dmda, heft and pheft for now) will use the performance models of all the implementations it was given, and pick the one that seems to be the fastest.

Enabling implementation according to capabilities
=================================================

Some implementations may not run on some devices. For instance, some CUDA devices do not support double floating point precision, and thus the kernel execution would just fail; or the device may not have enough shared memory for the implementation being used. The can_execute field of the struct starpu_codelet structure permits to express this. For instance:

     static int can_execute(unsigned workerid, struct starpu_task *task, unsigned nimpl)
     {
       const struct cudaDeviceProp *props;
       if (starpu_worker_get_type(workerid) == STARPU_CPU_WORKER)
         return 1;
       /* Cuda device */
       props = starpu_cuda_get_device_properties(workerid);
       if (props->major >= 2 || props->minor >= 3)
         /* At least compute capability 1.3, supports doubles */
         return 1;
       /* Old card, does not support doubles */
       return 0;
     }

     struct starpu_codelet cl = {
         .where = STARPU_CPU|STARPU_CUDA,
         .can_execute = can_execute,
         .cpu_funcs = { cpu_func, NULL },
         .cuda_funcs = { gpu_func, NULL }
         .nbuffers = 1,
         .modes = { STARPU_RW }
     };

This can be essential e.g. when running on a machine which mixes various models of CUDA devices, to take benefit from the new models without crashing on old models.

Note: the can_execute function is called by the scheduler each time it tries to match a task with a worker, and should thus be very fast. The starpu_cuda_get_device_properties provides a quick access to CUDA properties of CUDA devices to achieve such efficiency.

Another example is compiling CUDA code for various compute capabilities, resulting with two CUDA functions, e.g. scal_gpu_13 for compute capability 1.3, and scal_gpu_20 for compute capability 2.0. Both functions can be provided to StarPU by using cuda_funcs, and can_execute can then be used to rule out the scal_gpu_20 variant on a CUDA device which will not be able to execute it:

     static int can_execute(unsigned workerid, struct starpu_task *task, unsigned nimpl)
     {
       const struct cudaDeviceProp *props;
       if (starpu_worker_get_type(workerid) == STARPU_CPU_WORKER)
         return 1;
       /* Cuda device */
       if (nimpl == 0)
         /* Trying to execute the 1.3 capability variant, we assume it is ok in all cases.  */
         return 1;
       /* Trying to execute the 2.0 capability variant, check that the card can do it.  */
       props = starpu_cuda_get_device_properties(workerid);
       if (props->major >= 2 || props->minor >= 0)
         /* At least compute capability 2.0, can run it */
         return 1;
       /* Old card, does not support 2.0, will not be able to execute the 2.0 variant.  */
       return 0;
     }

     struct starpu_codelet cl = {
         .where = STARPU_CPU|STARPU_CUDA,
         .can_execute = can_execute,
         .cpu_funcs = { cpu_func, NULL },
         .cuda_funcs = { scal_gpu_13, scal_gpu_20, NULL },
         .nbuffers = 1,
         .modes = { STARPU_RW }
     };

Note: the most generic variant should be provided first, as some schedulers are not able to try the different variants.

Task and Worker Profiling
=========================

A full example showing how to use the profiling API is available in the StarPU sources in the directory examples/profiling/.

     struct starpu_task *task = starpu_task_create();
     task->cl = &cl;
     task->synchronous = 1;
     /* We will destroy the task structure by hand so that we can
      * query the profiling info before the task is destroyed. */
     task->destroy = 0;

     /* Submit and wait for completion (since synchronous was set to 1) */
     starpu_task_submit(task);

     /* The task is finished, get profiling information */
     struct starpu_task_profiling_info *info = task->profiling_info;

     /* How much time did it take before the task started ? */
     double delay += starpu_timing_timespec_delay_us(&info->submit_time, &info->start_time);

     /* How long was the task execution ? */
     double length += starpu_timing_timespec_delay_us(&info->start_time, &info->end_time);

     /* We don't need the task structure anymore */
     starpu_task_destroy(task);

     /* Display the occupancy of all workers during the test */
     int worker;
     for (worker = 0; worker < starpu_worker_get_count(); worker++)
     {
             struct starpu_worker_profiling_info worker_info;
             int ret = starpu_worker_get_profiling_info(worker, &worker_info);
             STARPU_ASSERT(!ret);

             double total_time = starpu_timing_timespec_to_us(&worker_info.total_time);
             double executing_time = starpu_timing_timespec_to_us(&worker_info.executing_time);
             double sleeping_time = starpu_timing_timespec_to_us(&worker_info.sleeping_time);

             float executing_ratio = 100.0*executing_time/total_time;
             float sleeping_ratio = 100.0*sleeping_time/total_time;

             char workername[128];
             starpu_worker_get_name(worker, workername, 128);
             fprintf(stderr, "Worker %s:\n", workername);
             fprintf(stderr, "\ttotal time: %.2lf ms\n", total_time*1e-3);
             fprintf(stderr, "\texec time: %.2lf ms (%.2f %%)\n", executing_time*1e-3,
                     executing_ratio);
             fprintf(stderr, "\tblocked time: %.2lf ms (%.2f %%)\n", sleeping_time*1e-3,
                     sleeping_ratio);
     }

Partitioning Data
=================

An existing piece of data can be partitioned in sub parts to be used by different tasks, for instance:

     int vector[NX];
     starpu_data_handle_t handle;

     /* Declare data to StarPU */
     starpu_vector_data_register(&handle, 0, (uintptr_t)vector, NX, sizeof(vector[0]));

     /* Partition the vector in PARTS sub-vectors */
     starpu_filter f =
     {
         .filter_func = starpu_block_filter_func_vector,
         .nchildren = PARTS
     };
     starpu_data_partition(handle, &f);

The task submission then uses starpu_data_get_sub_data to retrive the sub-handles to be passed as tasks parameters.

     /* Submit a task on each sub-vector */
     for (i=0; i<starpu_data_get_nb_children(handle); i++) {
         /* Get subdata number i (there is only 1 dimension) */
         starpu_data_handle_t sub_handle = starpu_data_get_sub_data(handle, 1, i);
         struct starpu_task *task = starpu_task_create();

         task->handles[0] = sub_handle;
         task->cl = &cl;
         task->synchronous = 1;
         task->cl_arg = &factor;
         task->cl_arg_size = sizeof(factor);

         starpu_task_submit(task);
     }

Partitioning can be applied several times, see examples/basic_examples/mult.c and examples/filters/.

Wherever the whole piece of data is already available, the partitioning will be done in-place, i.e. without allocating new buffers but just using pointers inside the existing copy. This is particularly important to be aware of when using OpenCL, where the kernel parameters are not pointers, but handles. The kernel thus needs to be also passed the offset within the OpenCL buffer:

     void opencl_func(void *buffers[], void *cl_arg)
     {
         cl_mem vector = (cl_mem) STARPU_VECTOR_GET_DEV_HANDLE(buffers[0]);
         unsigned offset = STARPU_BLOCK_GET_OFFSET(buffers[0]);

         ...
         clSetKernelArg(kernel, 0, sizeof(vector), &vector);
         clSetKernelArg(kernel, 1, sizeof(offset), &offset);
         ...
     }

And the kernel has to shift from the pointer passed by the OpenCL driver:

     __kernel void opencl_kernel(__global int *vector, unsigned offset)
     {
         block = (__global void *)block + offset;
         ...
     }

Performance model example
=========================

To achieve good scheduling, StarPU scheduling policies need to be able to estimate in advance the duration of a task. This is done by giving to codelets a performance model, by defining a starpu_perfmodel structure and providing its address in the model field of the struct starpu_codelet structure. The symbol and type fields of starpu_perfmodel are mandatory, to give a name to the model, and the type of the model, since there are several kinds of performance models.

  * Measured at runtime (STARPU_HISTORY_BASED model type). This assumes that for a given set of data input/output sizes, the performance will always be about the same. This is very true for regular kernels on GPUs for instance (<0.1% error), and just a bit less true on CPUs (~=1% error). This also assumes that there are few different sets of data input/output sizes. StarPU will then keep record of the average time of previous executions on the various processing units, and use it as an estimation. History is done per task size, by using a hash of the input and ouput sizes as an index. It will also save it in ~/.starpu/sampling/codelets for further executions, and can be observed by using the starpu_perfmodel_display command, or drawn by using the starpu_perfmodel_plot.  The models are indexed by machine name. To share the models between machines (e.g. for a homogeneous cluster), use export STARPU_HOSTNAME=some_global_name. Measurements are only done when using a task scheduler which makes use of it, such as heft or dmda.

     The following is a small code example.

     If e.g. the code is recompiled with other compilation options, or several variants of the code are used, the symbol string should be changed to reflect that, in order to recalibrate a new model from zero. The symbol string can even be constructed dynamically at execution time, as long as this is done before submitting any task using it.

          static struct starpu_perfmodel mult_perf_model = {
              .type = STARPU_HISTORY_BASED,
              .symbol = "mult_perf_model"
          };

          struct starpu_codelet cl = {
              .where = STARPU_CPU,
              .cpu_funcs = { cpu_mult, NULL },
              .nbuffers = 3,
              .modes = { STARPU_R, STARPU_R, STARPU_W },
              /* for the scheduling policy to be able to use performance models */
              .model = &mult_perf_model
          };

  * Measured at runtime and refined by regression (STARPU_*REGRESSION_BASED model type). This still assumes performance regularity, but can work with various data input sizes, by applying regression over observed execution times. STARPU_REGRESSION_BASED uses an a*n^b regression form, STARPU_NL_REGRESSION_BASED uses an a*n^b+c (more precise than STARPU_REGRESSION_BASED, but costs a lot more to compute). For instance, tests/perfmodels/regression_based.c uses a regression-based performance model for the memset operation. Of course, the application has to issue tasks with varying size so that the regression can be computed. StarPU will not trust the regression unless there is at least 10% difference between the minimum and maximum observed input size. For non-linear regression, since computing it is quite expensive, it is only done at termination of the application. This means that the first execution uses history-based performance model to perform scheduling.

  * Provided as an estimation from the application itself (STARPU_COMMON model type and cost_function field), see for instance examples/common/blas_model.h and examples/common/blas_model.c.

  * Provided explicitly by the application (STARPU_PER_ARCH model type): the .per_arch[arch][nimpl].cost_function fields have to be filled with pointers to functions which return the expected duration of the task in micro-seconds, one per architecture.

For the STARPU_HISTORY_BASED and STARPU_*REGRESSION_BASE, the total size of task data (both input and output) is used as an index by default. The size_base field of struct starpu_perfmodel however permits the application to override that, when for instance some of the data do not matter for task cost (e.g. mere reference table), or when using sparse structures (in which case it is the number of non-zeros which matter), or when there is some hidden parameter such as the number of iterations, etc.

How to use schedulers which can benefit from such performance model is explained in Task scheduling policy.

The same can be done for task power consumption estimation, by setting the power_model field the same way as the model field. Note: for now, the application has to give to the power consumption performance model a name which is different from the execution time performance model.

The application can request time estimations from the StarPU performance models by filling a task structure as usual without actually submitting it. The data handles can be created by calling starpu_data_register functions with a NULL pointer (and need to be unregistered as usual) and the desired data sizes. The starpu_task_expected_length and starpu_task_expected_power functions can then be called to get an estimation of the task duration on a given arch. starpu_task_destroy needs to be called to destroy the dummy task afterwards. See tests/perfmodels/regression_based.c for an example.

Theoretical lower bound on execution time
=========================================

For kernels with history-based performance models, StarPU can very easily provide a theoretical lower bound for the execution time of a whole set of tasks. See for instance examples/lu/lu_example.c: before submitting tasks, call starpu_bound_start, and after complete execution, call starpu_bound_stop. starpu_bound_print_lp or starpu_bound_print_mps can then be used to output a Linear Programming problem corresponding to the schedule of your tasks. Run it through lp_solve or any other linear programming solver, and that will give you a lower bound for the total execution time of your tasks. If StarPU was compiled with the glpk library installed, starpu_bound_compute can be used to solve it immediately and get the optimized minimum, in ms. Its integer parameter allows to decide whether integer resolution should be computed and returned too.

The deps parameter tells StarPU whether to take tasks and implicit data dependencies into account. It must be understood that the linear programming problem size is quadratic with the number of tasks and thus the time to solve it will be very long, it could be minutes for just a few dozen tasks. You should probably use lp_solve -timeout 1 test.pl -wmps test.mps to convert the problem to MPS format and then use a better solver, glpsol might be better than lp_solve for instance (the --pcost option may be useful), but sometimes doesn't manage to converge. cbc might look slower, but it is parallel. Be sure to try at least all the -B options of lp_solve. For instance, we often just use lp_solve -cc -B1 -Bb -Bg -Bp -Bf -Br -BG -Bd -Bs -BB -Bo -Bc -Bi , and the -gr option can also be quite useful.

Setting deps to 0 will only take into account the actual computations on processing units. It however still properly takes into account the varying performances of kernels and processing units, which is quite more accurate than just comparing StarPU performances with the fastest of the kernels being used.

The prio parameter tells StarPU whether to simulate taking into account the priorities as the StarPU scheduler would, i.e. schedule prioritized tasks before less prioritized tasks, to check to which extend this results to a less optimal solution. This increases even more computation time.

Note that for simplicity, all this however doesn't take into account data transfers, which are assumed to be completely overlapped.

Insert Task Utility
===================

StarPU provides the wrapper function starpu_insert_task to ease the creation and submission of tasks.

 -- Function: int starpu_insert_task (struct starpu_codelet *CL, ...)
     Create and submit a task corresponding to CL with the following arguments.  The argument list must be zero-terminated.

     The arguments following the codelets can be of the following types:

       * STARPU_R, STARPU_W, STARPU_RW, STARPU_SCRATCH, STARPU_REDUX an access mode followed by a data handle;

       * the specific values STARPU_VALUE, STARPU_CALLBACK, STARPU_CALLBACK_ARG, STARPU_CALLBACK_WITH_ARG, STARPU_PRIORITY, followed by the appropriated objects as defined below.

     Parameters to be passed to the codelet implementation are defined through the type STARPU_VALUE. The function starpu_codelet_unpack_args must be called within the codelet implementation to retrieve them.

 -- Macro: STARPU_VALUE
     this macro is used when calling starpu_insert_task, and must be followed by a pointer to a constant value and the size of the constant

 -- Macro: STARPU_CALLBACK
     this macro is used when calling starpu_insert_task, and must be followed by a pointer to a callback function

 -- Macro: STARPU_CALLBACK_ARG
     this macro is used when calling starpu_insert_task, and must be followed by a pointer to be given as an argument to the callback function

 -- Macro: STARPU_CALLBACK_WITH_ARG
     this macro is used when calling starpu_insert_task, and must be followed by two pointers: one to a callback function, and the other to be given as an argument to the callback function; this is equivalent to using both STARPU_CALLBACK and STARPU_CALLBACK_WITH_ARG

 -- Macro: STARPU_PRIORITY
     this macro is used when calling starpu_insert_task, and must be followed by a integer defining a priority level

 -- Function: void starpu_codelet_pack_args (char **ARG_BUFFER, size_t *ARG_BUFFER_SIZE, ...)
     Pack arguments of type STARPU_VALUE into a buffer which can be given to a codelet and later unpacked with the function starpu_codelet_unpack_args defined below.

 -- Function: void starpu_codelet_unpack_args (void *CL_ARG, ...)
     Retrieve the arguments of type STARPU_VALUE associated to a task automatically created using the function starpu_insert_task defined above.

Here the implementation of the codelet:

     void func_cpu(void *descr[], void *_args)
     {
             int *x0 = (int *)STARPU_VARIABLE_GET_PTR(descr[0]);
             float *x1 = (float *)STARPU_VARIABLE_GET_PTR(descr[1]);
             int ifactor;
             float ffactor;

             starpu_codelet_unpack_args(_args, &ifactor, &ffactor);
             *x0 = *x0 * ifactor;
             *x1 = *x1 * ffactor;
     }

     struct starpu_codelet mycodelet = {
             .where = STARPU_CPU,
             .cpu_funcs = { func_cpu, NULL },
             .nbuffers = 2,
             .modes = { STARPU_RW, STARPU_RW }
     };

And the call to the starpu_insert_task wrapper:

     starpu_insert_task(&mycodelet,
                        STARPU_VALUE, &ifactor, sizeof(ifactor),
                        STARPU_VALUE, &ffactor, sizeof(ffactor),
                        STARPU_RW, data_handles[0], STARPU_RW, data_handles[1],
                        0);

The call to starpu_insert_task is equivalent to the following code:

     struct starpu_task *task = starpu_task_create();
     task->cl = &mycodelet;
     task->handles[0] = data_handles[0];
     task->handles[1] = data_handles[1];
     char *arg_buffer;
     size_t arg_buffer_size;
     starpu_codelet_pack_args(&arg_buffer, &arg_buffer_size,
                         STARPU_VALUE, &ifactor, sizeof(ifactor),
                         STARPU_VALUE, &ffactor, sizeof(ffactor),
                         0);
     task->cl_arg = arg_buffer;
     task->cl_arg_size = arg_buffer_size;
     int ret = starpu_task_submit(task);

If some part of the task insertion depends on the value of some computation, the STARPU_DATA_ACQUIRE_CB macro can be very convenient. For instance, assuming that the index variable i was registered as handle i_handle:

     /* Compute which portion we will work on, e.g. pivot */
     starpu_insert_task(&which_index, STARPU_W, i_handle, 0);

     /* And submit the corresponding task */
     STARPU_DATA_ACQUIRE_CB(i_handle, STARPU_R, starpu_insert_task(&work, STARPU_RW, A_handle[i], 0));

The STARPU_DATA_ACQUIRE_CB macro submits an asynchronous request for acquiring data i for the main application, and will execute the code given as third parameter when it is acquired. In other words, as soon as the value of i computed by the which_index codelet can be read, the portion of code passed as third parameter of STARPU_DATA_ACQUIRE_CB will be executed, and is allowed to read from i to use it e.g. as an index. Note that this macro is only avaible when compiling StarPU with the compiler gcc.

Parallel Tasks
==============

StarPU can leverage existing parallel computation libraries by the means of parallel tasks. A parallel task is a task which gets worked on by a set of CPUs (called a parallel or combined worker) at the same time, by using an existing parallel CPU implementation of the computation to be achieved. This can also be useful to improve the load balance between slow CPUs and fast GPUs: since CPUs work collectively on a single task, the completion time of tasks on CPUs become comparable to the completion time on GPUs, thus relieving from granularity discrepancy concerns.

Two modes of execution exist to accomodate with existing usages.

Fork-mode parallel tasks
------------------------

In the Fork mode, StarPU will call the codelet function on one of the CPUs of the combined worker. The codelet function can use starpu_combined_worker_get_size() to get the number of threads it is allowed to start to achieve the computation. The CPU binding mask is already enforced, so that threads created by the function will inherit the mask, and thus execute where StarPU expected. For instance, using OpenMP (full source is available in examples/openmp/vector_scal.c):

     void scal_cpu_func(void *buffers[], void *_args)
     {
         unsigned i;
         float *factor = _args;
         struct starpu_vector_interface *vector = buffers[0];
         unsigned n = STARPU_VECTOR_GET_NX(vector);
         float *val = (float *)STARPU_VECTOR_GET_PTR(vector);

     #pragma omp parallel for num_threads(starpu_combined_worker_get_size())
         for (i = 0; i < n; i++)
             val[i] *= *factor;
     }

     static struct starpu_codelet cl =
     {
         .modes = { STARPU_RW },
         .where = STARPU_CPU,
         .type = STARPU_FORKJOIN,
         .max_parallelism = INT_MAX,
         .cpu_funcs = {scal_cpu_func, NULL},
         .nbuffers = 1,
     };

Other examples include for instance calling a BLAS parallel CPU implementation (see examples/mult/xgemm.c).

SPMD-mode parallel tasks
------------------------

In the SPMD mode, StarPU will call the codelet function on each CPU of the combined worker. The codelet function can use starpu_combined_worker_get_size() to get the total number of CPUs involved in the combined worker, and thus the number of calls that are made in parallel to the function, and starpu_combined_worker_get_rank() to get the rank of the current CPU within the combined worker. For instance:

     static void func(void *buffers[], void *args)
     {
         unsigned i;
         float *factor = _args;
         struct starpu_vector_interface *vector = buffers[0];
         unsigned n = STARPU_VECTOR_GET_NX(vector);
         float *val = (float *)STARPU_VECTOR_GET_PTR(vector);

         /* Compute slice to compute */
         unsigned m = starpu_combined_worker_get_size();
         unsigned j = starpu_combined_worker_get_rank();
         unsigned slice = (n+m-1)/m;

         for (i = j * slice; i < (j+1) * slice && i < n; i++)
             val[i] *= *factor;
     }

     static struct starpu_codelet cl =
     {
         .modes = { STARPU_RW },
         .where = STARP_CPU,
         .type = STARPU_SPMD,
         .max_parallelism = INT_MAX,
         .cpu_funcs = { func, NULL },
         .nbuffers = 1,
     }

Of course, this trivial example will not really benefit from parallel task execution, and was only meant to be simple to understand.  The benefit comes when the computation to be done is so that threads have to e.g. exchange intermediate results, or write to the data in a complex but safe way in the same buffer.

Parallel tasks performance
--------------------------

To benefit from parallel tasks, a parallel-task-aware StarPU scheduler has to be used. When exposed to codelets with a Fork or SPMD flag, the pheft (parallel-heft) and pgreedy (parallel greedy) schedulers will indeed also try to execute tasks with several CPUs. It will automatically try the various available combined worker sizes and thus be able to avoid choosing a large combined worker if the codelet does not actually scale so much.

Combined worker sizes
---------------------

By default, StarPU creates combined workers according to the architecture structure as detected by hwloc. It means that for each object of the hwloc topology (NUMA node, socket, cache, ...) a combined worker will be created. If some nodes of the hierarchy have a big arity (e.g. many cores in a socket without a hierarchy of shared caches), StarPU will create combined workers of intermediate sizes.

Concurrent parallel tasks
-------------------------

Unfortunately, many environments and librairies do not support concurrent calls.

For instance, most OpenMP implementations (including the main ones) do not support concurrent pragma omp parallel statements without nesting them in another pragma omp parallel statement, but StarPU does not yet support creating its CPU workers by using such pragma.

Other parallel libraries are also not safe when being invoked concurrently from different threads, due to the use of global variables in their sequential sections for instance.

The solution is then to use only a single combined worker, scoping all the CPUs.  This can be done by setting single_combined_worker to 1 in the starpu_conf structure, or setting the STARPU_SINGLE_COMBINED_WORKER environment variable to 1. StarPU will then use parallel tasks only over all the CPUs at the same time.

Debugging
=========

StarPU provides several tools to help debugging aplications. Execution traces can be generated and displayed graphically, see Generating traces. Some gdb helpers are also provided to show the whole StarPU state:

     (gdb) source tools/gdbinit
     (gdb) help starpu

The multiformat interface
=========================

It may be interesting to represent the same piece of data using two different data structures: one that would only be used on CPUs, and one that would only be used on GPUs. This can be done by using the multiformat interface. StarPU will be able to convert data from one data structure to the other when needed. Note that the heft scheduler is the only one optimized for this interface. The user must provide StarPU with conversion codelets:

     #define NX 1024
     struct point array_of_structs[NX];
     starpu_data_handle_t handle;

     /*
      * The conversion of a piece of data is itself a task, though it is created,
      * submitted and destroyed by StarPU internals and not by the user. Therefore,
      * we have to define two codelets.
      * Note that for now the conversion from the CPU format to the GPU format has to
      * be executed on the GPU, and the conversion from the GPU to the CPU has to be
      * executed on the CPU.
      */
     #ifdef STARPU_USE_OPENCL
     void cpu_to_opencl_opencl_func(void *buffers[], void *args);
     struct starpu_codelet cpu_to_opencl_cl = {
         .where = STARPU_OPENCL,
         .opencl_funcs = { cpu_to_opencl_opencl_func, NULL },
         .nbuffers = 1,
         .modes = { STARPU_RW }
     };

     void opencl_to_cpu_func(void *buffers[], void *args);
     struct starpu_codelet opencl_to_cpu_cl = {
         .where = STARPU_CPU,
         .cpu_funcs = { opencl_to_cpu_func, NULL },
         .nbuffers = 1,
         .modes = { STARPU_RW }
     };
     #endif

     struct starpu_multiformat_data_interface_ops format_ops = {
     #ifdef STARPU_USE_OPENCL
         .opencl_elemsize = 2 * sizeof(float),
         .cpu_to_opencl_cl = &cpu_to_opencl_cl,
         .opencl_to_cpu_cl = &opencl_to_cpu_cl,
     #endif
         .cpu_elemsize = 2 * sizeof(float),
         ...
     };
     starpu_multiformat_data_register(handle, 0, &array_of_structs, NX, &format_ops);

Kernels can be written almost as for any other interface. Note that STARPU_MULTIFORMAT_GET_PTR shall only be used for CPU kernels. CUDA kernels must use STARPU_MULTIFORMAT_GET_CUDA_PTR, and OpenCL kernels must use STARPU_MULTIFORMAT_GET_OPENCL_PTR. STARPU_MULTIFORMAT_GET_NX may be used in any kind of kernel.

     static void
     multiformat_scal_cpu_func(void *buffers[], void *args)
     {
         struct point *aos;
         unsigned int n;

         aos = STARPU_MULTIFORMAT_GET_PTR(buffers[0]);
         n = STARPU_MULTIFORMAT_GET_NX(buffers[0]);
         ...
     }

     extern "C" void multiformat_scal_cuda_func(void *buffers[], void *_args)
     {
         unsigned int n;
         struct struct_of_arrays *soa;

         soa = (struct struct_of_arrays *) STARPU_MULTIFORMAT_GET_CUDA_PTR(buffers[0]);
         n = STARPU_MULTIFORMAT_GET_NX(buffers[0]);

         ...
     }

A full example may be found in examples/basic_examples/multiformat.c.

On-GPU rendering
================

Graphical-oriented applications need to draw the result of their computations, typically on the very GPU where these happened. Technologies such as OpenGL/CUDA interoperability permit to let CUDA directly work on the OpenGL buffers, making them thus immediately ready for drawing, by mapping OpenGL buffer, textures or renderbuffer objects into CUDA. To achieve this with StarPU, it simply needs to be given the CUDA pointer at registration, for instance:

     for (workerid = 0; workerid < starpu_worker_get_count(); workerid++)
             if (starpu_worker_get_type(workerid) == STARPU_CUDA_WORKER)
                     break;

     cudaSetDevice(starpu_worker_get_devid(workerid));
     cudaGraphicsResourceGetMappedPointer((void**)&output, &num_bytes, resource);
     starpu_vector_data_register(&handle, starpu_worker_get_memory_node(workerid), output, num_bytes / sizeof(float4), sizeof(float4));

     starpu_insert_task(&cl, STARPU_RW, handle, 0);

     starpu_data_unregister(handle);

     cudaSetDevice(starpu_worker_get_devid(workerid));
     cudaGraphicsUnmapResources(1, &resource, 0);

     /* Now display it */

More examples
=============

More examples are available in the StarPU sources in the examples/ directory. Simple examples include:

incrementer/:
     Trivial incrementation test.

basic_examples/:
     Simple documented Hello world (as shown in Hello World), vector/scalar product (as shown in Vector Scaling on an Hybrid CPU/GPU Machine), matrix product examples (as shown in Performance model example), an example using the blocked matrix data interface, an example using the variable data interface, and an example using different formats on CPUs and GPUs.

matvecmult/:
     OpenCL example from NVidia, adapted to StarPU.

axpy/:
     AXPY CUBLAS operation adapted to StarPU.

fortran/:
     Example of Fortran bindings.

More advanced examples include:

filters/:
     Examples using filters, as shown in Partitioning Data.

lu/:
     LU matrix factorization, see for instance xlu_implicit.c

cholesky/:
     Cholesky matrix factorization, see for instance cholesky_implicit.c.