Add an example where it is possible to use StarPU with the runtime API.

This is *much* simpler than the driver API (it is actually likely that 99.9% of
code would use this API). It is probably much more efficient in terms of
register usage (in the driver API, we directly copy the data interface in the
arguments of the computational kernels, even though many parameters will remain
unused).

examples/Makefile.am

@@ -41,6 +41,10 @@ NVCC ?= nvcc
 	$(MKDIR_P) `dirname $@`
 	$(NVCC) -cubin $< -o $@ --compiler-options -fno-strict-aliasing  $(NVCCFLAGS)
 
+.cu.o:
+	$(NVCC) $< -c -o $@ --compiler-options -fno-strict-aliasing  $(NVCCFLAGS) -I${includedir}
+
+
 BUILT_SOURCES =					\
 	cuda/incrementer_cuda.cubin		\
 	cuda/spmv_cuda.cubin			
@@ -225,9 +229,23 @@ endif
 # Incrementer example #
 #######################
 
-check_PROGRAMS += incrementer/incrementer
 
-examplebin_PROGRAMS += incrementer/incrementer
+check_PROGRAMS +=				\
+	incrementer/incrementer			\
+	incrementer/incrementer_runtime
+
+examplebin_PROGRAMS +=				\
+	incrementer/incrementer			\
+	incrementer/incrementer_runtime
 
 incrementer_incrementer_SOURCES =		\
 	incrementer/incrementer.c
+
+if USE_CUDA
+incrementer_incrementer_runtime_SOURCES =	\
+	incrementer/incrementer_runtime.c	\
+	incrementer/incrementer_runtime_kernels.cu
+else
+incrementer_incrementer_runtime_SOURCES =	\
+	incrementer/incrementer_runtime.c
+endif

examples/incrementer/incrementer_runtime.c

@@ -0,0 +1,123 @@
+/*
+ * StarPU
+ * Copyright (C) INRIA 2008-2009 (see AUTHORS file)
+ *
+ * This program 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 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * This program 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 in COPYING.LGPL for more details.
+ */
+
+#include <string.h>
+#include <math.h>
+#include <sys/types.h>
+#include <pthread.h>
+#include <signal.h>
+
+/* for USE_CUDA */
+#include <starpu_config.h>
+
+#ifdef USE_CUDA
+#include <cuda.h>
+#endif
+
+#include <starpu.h>
+
+#define NITER	50000
+
+extern void cuda_codelet_host(float *tab);
+
+static starpu_data_handle my_float_state;
+
+static pthread_mutex_t mutex = PTHREAD_MUTEX_INITIALIZER;
+static pthread_cond_t cond = PTHREAD_COND_INITIALIZER;
+
+static float my_lovely_float[4] __attribute__ ((aligned (16))) = { 0.0f, 0.0f, 0.0f, 1664.0f}; 
+static unsigned i;
+
+void callback_func(void *argcb)
+{
+	unsigned cnt = STARPU_ATOMIC_ADD((unsigned *)argcb, 1);
+
+	if (cnt == NITER) 
+	{
+		pthread_mutex_lock(&mutex);
+		pthread_cond_signal(&cond);
+		pthread_mutex_unlock(&mutex);
+
+	}
+}
+
+void core_codelet(starpu_data_interface_t *buffers, __attribute__ ((unused)) void *_args)
+{
+	float *val = (float *)buffers[0].vector.ptr;
+
+	val[0] += 1.0f; val[1] += 1.0f;
+}
+
+#ifdef USE_CUDA
+void cuda_codelet(starpu_data_interface_t *buffers, __attribute__ ((unused)) void *_args)
+{
+	float *val = (float *)buffers[0].vector.ptr;
+
+	cuda_codelet_host(val);
+}
+
+#endif
+
+int main(__attribute__ ((unused)) int argc, __attribute__ ((unused)) char **argv)
+{
+	unsigned counter = 0;
+
+	starpu_init(NULL);
+
+	starpu_monitor_vector_data(&my_float_state, 0 /* home node */,
+			(uintptr_t)&my_lovely_float, 4, sizeof(float));
+
+	starpu_codelet cl =
+	{
+		/* CUBLAS stands for CUDA kernels controlled from the host */
+		.where = CORE|CUBLAS,
+		.core_func = core_codelet,
+#ifdef USE_CUDA
+		.cublas_func = &cuda_codelet,
+#endif
+		.nbuffers = 1
+	};
+
+	for (i = 0; i < NITER; i++)
+	{
+		struct starpu_task *task = starpu_task_create();
+		task->cl = &cl;
+		
+		task->callback_func = callback_func;
+		task->callback_arg = &counter;
+
+		task->buffers[0].state = my_float_state;
+		task->buffers[0].mode = RW;
+
+		starpu_submit_task(task);
+	}
+
+	pthread_mutex_lock(&mutex);
+	pthread_cond_wait(&cond, &mutex);
+	pthread_mutex_unlock(&mutex);
+
+	starpu_sync_data_with_mem(my_float_state);
+	
+	printf("array -> %f, %f, %f\n", my_lovely_float[0], 
+			my_lovely_float[1], my_lovely_float[2]);
+	
+	if (my_lovely_float[0] != my_lovely_float[1] + my_lovely_float[2])
+		return 1;
+	
+	starpu_shutdown();
+
+	return 0;
+}

examples/incrementer/incrementer_runtime_kernels.cu

@@ -0,0 +1,28 @@
+/*
+ * StarPU
+ * Copyright (C) INRIA 2008-2009 (see AUTHORS file)
+ *
+ * This program 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 2.1 of the License, or (at
+ * your option) any later version.
+ *
+ * This program 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 in COPYING.LGPL for more details.
+ */
+
+extern "C" __global__ void cuda_incrementer(float * tab)
+{
+	tab[0] = tab[0] + 1.0;
+	tab[2] = tab[2] + 1.0;
+	
+	return;
+}
+
+extern "C" void cuda_codelet_host(float *tab)
+{
+	cuda_incrementer<<<1,1>>>(tab);
+}