StarPU Handbook

Basic Examples

# Hello World

This section shows how to implement a simple program that submits a task to StarPU.

## Required Headers

The header starpu.h should be included in any code using StarPU.

#include <starpu.h>
starpu.h

## Defining A Codelet

A codelet is a structure that represents a computational kernel. Such a codelet may contain an implementation of the same kernel on different architectures (e.g. CUDA, x86, ...). For compatibility, make sure that the whole structure is properly initialized to zero, either by using the function starpu_codelet_init(), or by letting the compiler implicitly do it as examplified below.

The field starpu_codelet::nbuffers specifies the number of data buffers that are manipulated by the codelet: here the codelet does not access or modify any data that is controlled by our data management library.

We create a codelet which may only be executed on CPUs. When a CPU core will execute a codelet, it will call the function `cpu_func`, which must have the following prototype:

void (*cpu_func)(void *buffers[], void *cl_arg);

In this example, we can ignore the first argument of this function which gives a description of the input and output buffers (e.g. the size and the location of the matrices) since there is none. We also ignore the second argument which is a pointer to optional arguments for the codelet.

void cpu_func(void *buffers[], void *cl_arg)
{
printf("Hello world\n");
}
struct starpu_codelet cl =
{
.cpu_funcs = { cpu_func },
.nbuffers = 0
};
starpu_codelet::cpu_funcs
starpu_cpu_func_t cpu_funcs[STARPU_MAXIMPLEMENTATIONS]
Definition: starpu_task.h:326
starpu_codelet::cpu_func
starpu_cpu_func_t cpu_func
Definition: starpu_task.h:295
starpu_codelet
Definition: starpu_task.h:250

## Submitting A Task

Before submitting any tasks to StarPU, starpu_init() must be called. The `NULL` argument specifies that we use the default configuration. Tasks can then be submitted until the termination of StarPU – done by a call to starpu_shutdown().

In the example below, a task structure is allocated by a call to starpu_task_create(). This function allocates and fills the task structure with its default settings, it does not submit the task to StarPU.

The field starpu_task::cl is a pointer to the codelet which the task will execute: in other words, the codelet structure describes which computational kernel should be offloaded on the different architectures, and the task structure is a wrapper containing a codelet and the piece of data on which the codelet should operate.

If the field starpu_task::synchronous is non-zero, task submission will be synchronous: the function starpu_task_submit() will not return until the task has been executed. Note that the function starpu_shutdown() does not guarantee that asynchronous tasks have been executed before it returns, starpu_task_wait_for_all() can be used to this effect, or data can be unregistered (starpu_data_unregister()), which will implicitly wait for all the tasks scheduled to work on it, unless explicitly disabled thanks to starpu_data_set_default_sequential_consistency_flag() or starpu_data_set_sequential_consistency_flag().

int main(int argc, char **argv)
{
/* initialize StarPU */
starpu_init(NULL);
struct starpu_task *task = starpu_task_create();
task->cl = &cl; /* Pointer to the codelet defined above */
/* starpu_task_submit will be a blocking call. If unset,
starpu_task_wait() needs to be called after submitting the task. */
task->synchronous = 1;
/* submit the task to StarPU */
starpu_task_submit(task);
/* terminate StarPU */
starpu_shutdown();
return 0;
}
starpu_task::synchronous
unsigned synchronous
Definition: starpu_task.h:907
starpu_task::cl
struct starpu_codelet * cl
Definition: starpu_task.h:584
starpu_task_create
struct starpu_task * starpu_task_create(void) STARPU_ATTRIBUTE_MALLOC
starpu_task_submit
int starpu_task_submit(struct starpu_task *task) STARPU_WARN_UNUSED_RESULT
starpu_task
Definition: starpu_task.h:567
starpu_shutdown
void starpu_shutdown(void)
starpu_init
int starpu_init(struct starpu_conf *conf) STARPU_WARN_UNUSED_RESULT

## Execution Of Hello World

    $ make hello_world
    cc $(pkg-config --cflags starpu-1.3) hello_world.c -o hello_world $(pkg-config --libs starpu-1.3)
    $ ./hello_world
    Hello world

## Passing Arguments To The Codelet

The optional field starpu_task::cl_arg field is a pointer to a buffer (of size starpu_task::cl_arg_size) with some parameters for the kernel described by the codelet. For instance, if a codelet implements a computational kernel that multiplies its input vector by a constant, the constant could be specified by the means of this buffer, instead of registering it as a StarPU data. It must however be noted that StarPU avoids making copy whenever possible and rather passes the pointer as such, so the buffer which is pointed at must be kept allocated until the task terminates, and if several tasks are submitted with various parameters, each of them must be given a pointer to their own buffer.

struct params
{
int i;
float f;
};
void cpu_func(void *buffers[], void *cl_arg)
{
struct params *params = cl_arg;
printf("Hello world (params = {%i, %f} )\n", params->i, params->f);
}

As said before, the field starpu_codelet::nbuffers specifies the number of data buffers which are manipulated by the codelet. It does not count the argument — the parameter `cl_arg` of the function `cpu_func` — since it is not managed by our data management library, but just contains trivial parameters.

Be aware that this may be a pointer to a copy of the actual buffer, and not the pointer given by the programmer: if the codelet modifies this buffer, there is no guarantee that the initial buffer will be modified as well: this for instance implies that the buffer cannot be used as a synchronization medium. If synchronization is needed, data has to be registered to StarPU, see Vector Scaling.

int main(int argc, char **argv)
{
/* initialize StarPU */
starpu_init(NULL);
struct starpu_task *task = starpu_task_create();
task->cl = &cl; /* Pointer to the codelet defined above */
struct params params = { 1, 2.0f };
task->cl_arg = &params;
task->cl_arg_size = sizeof(params);
/* starpu_task_submit will be a blocking call */
task->synchronous = 1;
/* submit the task to StarPU */
starpu_task_submit(task);
/* terminate StarPU */
starpu_shutdown();
return 0;
}
starpu_task::cl_arg
void * cl_arg
Definition: starpu_task.h:709
starpu_task::cl_arg_size
size_t cl_arg_size
Definition: starpu_task.h:726

    $ make hello_world
    cc $(pkg-config --cflags starpu-1.3) hello_world.c -o hello_world $(pkg-config --libs starpu-1.3)
    $ ./hello_world
    Hello world (params = {1, 2.000000} )

## Defining A Callback

Once a task has been executed, an optional callback function starpu_task::callback_func is called when defined. While the computational kernel could be offloaded on various architectures, the callback function is always executed on a CPU. The pointer starpu_task::callback_arg is passed as an argument of the callback function. The prototype of a callback function must be:

void (*callback_function)(void *);
void callback_func(void *callback_arg)
{
printf("Callback function (arg %x)\n", callback_arg);
}
int main(int argc, char **argv)
{
/* initialize StarPU */
starpu_init(NULL);
struct starpu_task *task = starpu_task_create();
task->cl = &cl; /* Pointer to the codelet defined above */
task->callback_func = callback_func;
task->callback_arg = 0x42;
/* starpu_task_submit will be a blocking call */
task->synchronous = 1;
/* submit the task to StarPU */
starpu_task_submit(task);
/* terminate StarPU */
starpu_shutdown();
return 0;
}
starpu_task::callback_func
void(* callback_func)(void *)
Definition: starpu_task.h:744
starpu_task::callback_arg
void * callback_arg
Definition: starpu_task.h:757

    $ make hello_world
    cc $(pkg-config --cflags starpu-1.3) hello_world.c -o hello_world $(pkg-config --libs starpu-1.3)
    $ ./hello_world
    Hello world
    Callback function (arg 42)

## Where To Execute A Codelet

struct starpu_codelet cl =
{
.where = STARPU_CPU,
.cpu_funcs = { cpu_func },
.cpu_funcs_name = { "cpu_func" },
.nbuffers = 0
};
starpu_codelet::where
uint32_t where
Definition: starpu_task.h:265
STARPU_CPU
#define STARPU_CPU
Definition: starpu_task.h:55

We create a codelet which may only be executed on the CPUs. The optional field starpu_codelet::where is a bitmask which defines where the codelet may be executed. Here, the value STARPU_CPU means that only CPUs can execute this codelet. When the optional field starpu_codelet::where is unset, its value is automatically set based on the availability of the different fields `XXX_funcs`.

TODO: explain starpu_codelet::cpu_funcs_name

# Vector Scaling

The previous example has shown how to submit tasks. In this section, we show how StarPU tasks can manipulate data.

The full source code for this example is given in Full source code for the ’Scaling a Vector’ example.

## Source Code of Vector Scaling

Programmers can describe the data layout of their application so that StarPU is responsible for enforcing data coherency and availability across the machine. Instead of handling complex (and non-portable) mechanisms to perform data movements, programmers only declare which piece of data is accessed and/or modified by a task, and StarPU makes sure that when a computational kernel starts somewhere (e.g. on a GPU), its data are available locally.

Before submitting those tasks, the programmer first needs to declare the different pieces of data to StarPU using the functions `starpu_*_data_register`. To ease the development of applications for StarPU, it is possible to describe multiple types of data layout. A type of data layout is called an interface. There are different predefined interfaces available in StarPU: here we will consider the vector interface.

The following lines show how to declare an array of `NX` elements of type `float` using the vector interface:

float vector[NX];
starpu_data_handle_t vector_handle;
starpu_vector_data_register(&vector_handle, STARPU_MAIN_RAM, (uintptr_t)vector, NX, sizeof(vector[0]));
STARPU_MAIN_RAM
#define STARPU_MAIN_RAM
Definition: starpu_task.h:119
starpu_vector_data_register
void starpu_vector_data_register(starpu_data_handle_t *handle, int home_node, uintptr_t ptr, uint32_t nx, size_t elemsize)
starpu_data_handle_t
struct _starpu_data_state * starpu_data_handle_t
Definition: starpu_data.h:44

The first argument, called the data handle, is an opaque pointer which designates the array within StarPU. This is also the structure which is used to describe which data is used by a task. The second argument is the node number where the data originally resides. Here it is STARPU_MAIN_RAM since the array `vector` is in the main memory. Then comes the pointer `vector` where the data can be found in main memory, the number of elements in the vector and the size of each element. The following shows how to construct a StarPU task that will manipulate the vector and a constant factor.

float factor = 3.14;
struct starpu_task *task = starpu_task_create();
task->cl = &cl; /* Pointer to the codelet defined below */
task->handles[0] = vector_handle; /* First parameter of the codelet */
task->cl_arg = &factor;
task->cl_arg_size = sizeof(factor);
task->synchronous = 1;
starpu_task_submit(task);
starpu_task::handles
starpu_data_handle_t handles[STARPU_NMAXBUFS]
Definition: starpu_task.h:661

Since the factor is a mere constant float value parameter, it does not need a preliminary registration, and can just be passed through the pointer starpu_task::cl_arg like in the previous example. The vector parameter is described by its handle. starpu_task::handles should be set with the handles of the data, the access modes for the data are defined in the field starpu_codelet::modes (STARPU_R for read-only, STARPU_W for write-only and STARPU_RW for read and write access).

The definition of the codelet can be written as follows:

void scal_cpu_func(void *buffers[], void *cl_arg)
{
unsigned i;
float *factor = cl_arg;
/* length of the vector */
unsigned n = STARPU_VECTOR_GET_NX(buffers[0]);
/* CPU copy of the vector pointer */
float *val = (float *)STARPU_VECTOR_GET_PTR(buffers[0]);
for (i = 0; i < n; i++)
val[i] *= *factor;
}
struct starpu_codelet cl =
{
.cpu_funcs = { scal_cpu_func },
.cpu_funcs_name = { "scal_cpu_func" },
.nbuffers = 1,
.modes = { STARPU_RW }
};
STARPU_VECTOR_GET_NX
#define STARPU_VECTOR_GET_NX(interface)
Definition: starpu_data_interfaces.h:1345
STARPU_VECTOR_GET_PTR
#define STARPU_VECTOR_GET_PTR(interface)
Definition: starpu_data_interfaces.h:1329
STARPU_RW
@ STARPU_RW
Definition: starpu_data.h:59

The first argument is an array that gives a description of all the buffers passed in the array starpu_task::handles. The size of this array is given by the field starpu_codelet::nbuffers. For the sake of genericity, this array contains pointers to the different interfaces describing each buffer. In the case of the vector interface, the location of the vector (resp. its length) is accessible in the starpu_vector_interface::ptr (resp. starpu_vector_interface::nx) of this interface. Since the vector is accessed in a read-write fashion, any modification will automatically affect future accesses to this vector made by other tasks.

The second argument of the function `scal_cpu_func` contains a pointer to the parameters of the codelet (given in starpu_task::cl_arg), so that we read the constant factor from this pointer.

## Execution of Vector Scaling

    $ make vector_scal
    cc $(pkg-config --cflags starpu-1.3) vector_scal.c -o vector_scal $(pkg-config --libs starpu-1.3)
    $ ./vector_scal
    0.000000 3.000000 6.000000 9.000000 12.000000

# Vector Scaling on an Hybrid CPU/GPU Machine

Contrary to the previous examples, the task submitted in this example may not only be executed by the CPUs, but also by a CUDA device.

## Definition of the CUDA Kernel

The CUDA implementation can be written as follows. It needs to be compiled with a CUDA compiler such as nvcc, the NVIDIA CUDA compiler driver. It must be noted that the vector pointer returned by STARPU_VECTOR_GET_PTR is here a pointer in GPU memory, so that it can be passed as such to the kernel call `vector_mult_cuda`.

#include <starpu.h>
static __global__ void vector_mult_cuda(unsigned n, float *val, float factor)
{
unsigned i = blockIdx.x*blockDim.x + threadIdx.x;
if (i < n)
val[i] *= factor;
}
extern "C" void scal_cuda_func(void *buffers[], void *_args)
{
float *factor = (float *)_args;
/* length of the vector */
unsigned n = STARPU_VECTOR_GET_NX(buffers[0]);
/* local copy of the vector pointer */
float *val = (float *)STARPU_VECTOR_GET_PTR(buffers[0]);
unsigned threads_per_block = 64;
unsigned nblocks = (n + threads_per_block-1) / threads_per_block;
vector_mult_cuda<<<nblocks,threads_per_block, 0, starpu_cuda_get_local_stream()>>>(n, val, *factor);
cudaError_t status = cudaGetLastError();
if (status != cudaSuccess) STARPU_CUDA_REPORT_ERROR(status);
cudaStreamSynchronize(starpu_cuda_get_local_stream());
}
starpu_cuda_get_local_stream
cudaStream_t starpu_cuda_get_local_stream(void)
STARPU_CUDA_REPORT_ERROR
#define STARPU_CUDA_REPORT_ERROR(status)
Definition: starpu_cuda.h:59

## Definition of the OpenCL Kernel

The OpenCL implementation can be written as follows. StarPU provides tools to compile a OpenCL kernel stored in a file.

__kernel void vector_mult_opencl(int nx, __global float* val, float factor)
{
const int i = get_global_id(0);
if (i < nx)
{
val[i] *= factor;
}
}

Contrary to CUDA and CPU, STARPU_VECTOR_GET_DEV_HANDLE has to be used, which returns a `cl_mem` (which is not a device pointer, but an OpenCL handle), which can be passed as such to the OpenCL kernel. The difference is important when using partitioning, see Partitioning Data.

#include <starpu.h>
extern struct starpu_opencl_program programs;
void scal_opencl_func(void *buffers[], void *_args)
{
float *factor = _args;
int id, devid, err; /* OpenCL specific code */
cl_kernel kernel; /* OpenCL specific code */
cl_command_queue queue; /* OpenCL specific code */
cl_event event; /* OpenCL specific code */
/* length of the vector */
unsigned n = STARPU_VECTOR_GET_NX(buffers[0]);
/* OpenCL copy of the vector pointer */
cl_mem val = (cl_mem)STARPU_VECTOR_GET_DEV_HANDLE(buffers[0]);
{ /* OpenCL specific code */
id = starpu_worker_get_id();
devid = starpu_worker_get_devid(id);
err = starpu_opencl_load_kernel(&kernel, &queue, &programs,
"vector_mult_opencl", /* Name of the codelet */
devid);
if (err != CL_SUCCESS) STARPU_OPENCL_REPORT_ERROR(err);
err = clSetKernelArg(kernel, 0, sizeof(n), &n);
err |= clSetKernelArg(kernel, 1, sizeof(val), &val);
err |= clSetKernelArg(kernel, 2, sizeof(*factor), factor);
if (err) STARPU_OPENCL_REPORT_ERROR(err);
}
{ /* OpenCL specific code */
size_t global=n;
size_t local;
size_t s;
cl_device_id device;
starpu_opencl_get_device(devid, &device);
err = clGetKernelWorkGroupInfo (kernel, device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(local), &local, &s);
if (err != CL_SUCCESS) STARPU_OPENCL_REPORT_ERROR(err);
if (local > global) local=global;
else global = (global + local-1) / local * local;
err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, &local, 0, NULL, &event);
if (err != CL_SUCCESS) STARPU_OPENCL_REPORT_ERROR(err);
}
{ /* OpenCL specific code */
clFinish(queue);
starpu_opencl_collect_stats(event);
clReleaseEvent(event);
starpu_opencl_release_kernel(kernel);
}
}
STARPU_VECTOR_GET_DEV_HANDLE
#define STARPU_VECTOR_GET_DEV_HANDLE(interface)
Definition: starpu_data_interfaces.h:1335
starpu_opencl_program::programs
cl_program programs[STARPU_MAXOPENCLDEVS]
Definition: starpu_opencl.h:49
starpu_opencl_load_kernel
int starpu_opencl_load_kernel(cl_kernel *kernel, cl_command_queue *queue, struct starpu_opencl_program *opencl_programs, const char *kernel_name, int devid)
STARPU_OPENCL_REPORT_ERROR
#define STARPU_OPENCL_REPORT_ERROR(status)
Definition: starpu_opencl.h:259
starpu_opencl_release_kernel
int starpu_opencl_release_kernel(cl_kernel kernel)
starpu_opencl_get_device
void starpu_opencl_get_device(int devid, cl_device_id *device)
starpu_opencl_collect_stats
int starpu_opencl_collect_stats(cl_event event)
starpu_opencl_program
Definition: starpu_opencl.h:47
starpu_worker_get_devid
int starpu_worker_get_devid(int id)
starpu_worker_get_id
int starpu_worker_get_id(void)

## Definition of the Main Code

The CPU implementation is the same as in the previous section.

Here is the source of the main application. You can notice that the fields starpu_codelet::cuda_funcs and starpu_codelet::opencl_funcs are set to define the pointers to the CUDA and OpenCL implementations of the task.

/*
* This example demonstrates how to use StarPU to scale an array by a factor.
* It shows how to manipulate data with StarPU's data management library.
* 1- how to declare a piece of data to StarPU (starpu_vector_data_register)
* 2- how to describe which data are accessed by a task (task->handles[0])
* 3- how a kernel can manipulate the data (buffers[0].vector.ptr)
*/
#include <starpu.h>
#define NX 2048
extern void scal_cpu_func(void *buffers[], void *_args);
extern void scal_sse_func(void *buffers[], void *_args);
extern void scal_cuda_func(void *buffers[], void *_args);
extern void scal_opencl_func(void *buffers[], void *_args);
static struct starpu_codelet cl =
{
.where = STARPU_CPU | STARPU_CUDA | STARPU_OPENCL,
/* CPU implementation of the codelet */
.cpu_funcs = { scal_cpu_func, scal_sse_func },
.cpu_funcs_name = { "scal_cpu_func", "scal_sse_func" },
#ifdef STARPU_USE_CUDA
/* CUDA implementation of the codelet */
.cuda_funcs = { scal_cuda_func },
#endif
#ifdef STARPU_USE_OPENCL
/* OpenCL implementation of the codelet */
.opencl_funcs = { scal_opencl_func },
#endif
.nbuffers = 1,
.modes = { STARPU_RW }
};
#ifdef STARPU_USE_OPENCL
struct starpu_opencl_program programs;
#endif
int main(int argc, char **argv)
{
/* We consider a vector of float that is initialized just as any of C
* data */
float vector[NX];
unsigned i;
for (i = 0; i < NX; i++)
vector[i] = 1.0f;
fprintf(stderr, "BEFORE: First element was %f\n", vector[0]);
/* Initialize StarPU with default configuration */
starpu_init(NULL);
#ifdef STARPU_USE_OPENCL
starpu_opencl_load_opencl_from_file("examples/basic_examples/vector_scal_opencl_kernel.cl", &programs, NULL);
#endif
/* Tell StaPU to associate the "vector" vector with the "vector_handle"
* identifier. When a task needs to access a piece of data, it should
* refer to the handle that is associated to it.
* In the case of the "vector" data interface:
* - the first argument of the registration method is a pointer to the
* handle that should describe the data
* - the second argument is the memory node where the data (ie. "vector")
* resides initially: STARPU_MAIN_RAM stands for an address in main memory, as
* opposed to an adress on a GPU for instance.
* - the third argument is the adress of the vector in RAM
* - the fourth argument is the number of elements in the vector
* - the fifth argument is the size of each element.
*/
starpu_data_handle_t vector_handle;
starpu_vector_data_register(&vector_handle, STARPU_MAIN_RAM, (uintptr_t)vector, NX, sizeof(vector[0]));
float factor = 3.14;
/* create a synchronous task: any call to starpu_task_submit will block
* until it is terminated */
struct starpu_task *task = starpu_task_create();
task->synchronous = 1;
task->cl = &cl;
/* the codelet manipulates one buffer in RW mode */
task->handles[0] = vector_handle;
/* an argument is passed to the codelet, beware that this is a
* READ-ONLY buffer and that the codelet may be given a pointer to a
* COPY of the argument */
task->cl_arg = &factor;
task->cl_arg_size = sizeof(factor);
/* execute the task on any eligible computational ressource */
starpu_task_submit(task);
/* StarPU does not need to manipulate the array anymore so we can stop
* monitoring it */
starpu_data_unregister(vector_handle);
#ifdef STARPU_USE_OPENCL
starpu_opencl_unload_opencl(&programs);
#endif
/* terminate StarPU, no task can be submitted after */
starpu_shutdown();
fprintf(stderr, "AFTER First element is %f\n", vector[0]);
return 0;
}
STARPU_CUDA
#define STARPU_CUDA
Definition: starpu_task.h:62
STARPU_OPENCL
#define STARPU_OPENCL
Definition: starpu_task.h:69
starpu_data_unregister
void starpu_data_unregister(starpu_data_handle_t handle)
starpu_opencl_load_opencl_from_file
int starpu_opencl_load_opencl_from_file(const char *source_file_name, struct starpu_opencl_program *opencl_programs, const char *build_options)
starpu_opencl_unload_opencl
int starpu_opencl_unload_opencl(struct starpu_opencl_program *opencl_programs)

## Execution of Hybrid Vector Scaling

The Makefile given at the beginning of the section must be extended to give the rules to compile the CUDA source code. Note that the source file of the OpenCL kernel does not need to be compiled now, it will be compiled at run-time when calling the function starpu_opencl_load_opencl_from_file().

    CFLAGS  += $(shell pkg-config --cflags starpu-1.3)
    LDLIBS  += $(shell pkg-config --libs starpu-1.3)
    CC       = gcc

    vector_scal: vector_scal.o vector_scal_cpu.o vector_scal_cuda.o vector_scal_opencl.o

    %.o: %.cu
           nvcc $(CFLAGS) $< -c $@

    clean:
           rm -f vector_scal *.o

    $ make

and to execute it, with the default configuration:

    $ ./vector_scal
    0.000000 3.000000 6.000000 9.000000 12.000000

or for example, by disabling CPU devices:

    $ STARPU_NCPU=0 ./vector_scal
    0.000000 3.000000 6.000000 9.000000 12.000000

or by disabling CUDA devices (which may permit to enable the use of OpenCL, see Enabling OpenCL) :

    $ STARPU_NCUDA=0 ./vector_scal
    0.000000 3.000000 6.000000 9.000000 12.000000