examples/reductions/dot_product: add a cuda codelet for the reduction

examples/Makefile.am

@@ -629,6 +629,10 @@ examplebin_PROGRAMS +=				\
 
 reductions_dot_product_SOURCES =		\
 	reductions/dot_product.c
+if STARPU_USE_CUDA
+reductions_dot_product_SOURCES +=		\
+	reductions/dot_product_kernels.cu
+endif
 
 #####################
 # Min/Max reduction #

examples/reductions/dot_product.c

@@ -77,9 +77,16 @@ void redux_cpu_func(void *descr[], void *cl_arg)
 	*dota = *dota + *dotb;
 }
 
+#ifdef STARPU_USE_CUDA
+extern void redux_cuda_func(void *descr[], void *_args);
+#endif
+
 static struct starpu_codelet_t redux_codelet = {
-	.where = STARPU_CPU,
+	.where = STARPU_CPU|STARPU_CUDA,
 	.cpu_func = redux_cpu_func,
+#ifdef STARPU_USE_CUDA
+	.cuda_func = redux_cuda_func,
+#endif
 	.nbuffers = 2
 };
 
@@ -148,8 +155,6 @@ static struct starpu_codelet_t dot_codelet = {
  *	Tasks initialization
  */
 
-extern void starpu_data_end_reduction_mode(starpu_data_handle handle);
-
 int main(int argc, char **argv)
 {
 	starpu_init(NULL);

examples/reductions/dot_product_kernels.cu

@@ -0,0 +1,36 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2011  Centre National de la Recherche Scientifique
+ *
+ * StarPU 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.
+ *
+ * StarPU 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 <starpu.h>
+#include <starpu_cuda.h>
+
+#define DOT_TYPE double
+
+static __global__ void cuda_redux(DOT_TYPE *dota, DOT_TYPE *dotb)
+{
+	*dota = *dota + *dotb;
+	return;
+}
+
+extern "C" void redux_cuda_func(void *descr[], void *_args)
+{
+	(void)_args;
+	DOT_TYPE *dota = (DOT_TYPE *)STARPU_VARIABLE_GET_PTR(descr[0]);
+	DOT_TYPE *dotb = (DOT_TYPE *)STARPU_VARIABLE_GET_PTR(descr[1]);
+
+	cuda_redux<<<1,1, 0, starpu_cuda_get_local_stream()>>>(dota, dotb);
+	cudaStreamSynchronize(starpu_cuda_get_local_stream());
+}