Make examples/axpy/axpy work with OpenCL.

examples/Makefile.am

@@ -31,6 +31,7 @@ nobase_STARPU_OPENCL_DATA_DATA =
 endif
 
 EXTRA_DIST = 					\
+	axpy/axpy_opencl_kernel.cl \
 	basic_examples/vector_scal_opencl_kernel.cl \
 	basic_examples/multiformat_opencl_kernel.cl  \
 	basic_examples/multiformat_conversion_codelets_opencl_kernel.cl \
@@ -102,6 +103,7 @@ examplebindir = $(libdir)/starpu/examples/
 examplebin_PROGRAMS =
 
 noinst_HEADERS = 				\
+	axpy/axpy.h                             \
 	cg/cg.h					\
 	heat/lu_kernels_model.h			\
 	heat/dw_sparse_cg.h			\
@@ -439,6 +441,12 @@ if !NO_BLAS_LIB
 axpy_axpy_SOURCES =				\
 	axpy/axpy.c				\
 	common/blas.c
+if STARPU_USE_OPENCL
+axpy_axpy_SOURCES+=\
+	axpy/axpy_opencl.c
+nobase_STARPU_OPENCL_DATA_DATA+=
+	axpy/axpy_opencl_kernel.cl
+endif
 
 axpy_axpy_LDADD =				\
 	$(STARPU_BLAS_LDFLAGS)

examples/axpy/axpy.c

@@ -27,8 +27,12 @@
 #ifdef STARPU_USE_CUDA
 #include <cublas.h>
 #endif
+#ifdef STARPU_USE_OPENCL
+#include <starpu_opencl.h>
+#endif
+
+#include "axpy.h"
 
-#define TYPE	float
 #define AXPY	SAXPY
 #define CUBLASAXPY	cublasSaxpy
 
@@ -73,18 +77,19 @@ void axpy_gpu(void *descr[], __attribute__((unused)) void *arg)
 }
 #endif
 
-static struct starpu_codelet axpy_cl =
-{
-        .where =
-#ifdef STARPU_USE_CUDA
-                STARPU_CUDA|
+#ifdef STARPU_USE_OPENCL
+extern void axpy_opencl(void *buffers[], void *args);
 #endif
-                STARPU_CPU,
 
+static struct starpu_codelet axpy_cl =
+{
 	.cpu_funcs = {axpy_cpu, NULL},
 #ifdef STARPU_USE_CUDA
 	.cuda_funcs = {axpy_gpu, NULL},
 #endif
+#ifdef STARPU_USE_OPENCL
+	.opencl_funcs = {axpy_opencl, NULL},
+#endif
 	.nbuffers = 2,
 	.modes = {STARPU_R, STARPU_RW}
 };
@@ -103,6 +108,10 @@ check(void)
 	return EXIT_SUCCESS;
 }
 
+#ifdef STARPU_USE_OPENCL
+struct starpu_opencl_program opencl_program;
+#endif
+
 int main(int argc, char **argv)
 {
 	int ret;
@@ -113,6 +122,12 @@ int main(int argc, char **argv)
 		return 77;
 	STARPU_CHECK_RETURN_VALUE(ret, "starpu_init");
 
+#ifdef STARPU_USE_OPENCL
+	ret = starpu_opencl_load_opencl_from_file("examples/axpy/axpy_opencl_kernel.cl",
+						  &opencl_program, NULL);
+	STARPU_CHECK_RETURN_VALUE(ret, "starpu_opencl_load_opencl_from_file");
+#endif
+
 	starpu_helper_cublas_init();
 
 	/* This is equivalent to
@@ -197,6 +212,9 @@ enodev:
 	starpu_free((void *)vec_x);
 	starpu_free((void *)vec_y);
 
+#ifdef STARPU_USE_OPENCL
+        starpu_opencl_unload_opencl(&opencl_program);
+#endif
 	/* Stop StarPU */
 	starpu_shutdown();
 

examples/axpy/axpy.h

@@ -0,0 +1,7 @@
+#ifndef AXPY_H__
+#define AXPY_H__
+
+#define TYPE float
+
+#endif /* AXPY_H__ */
+

examples/axpy/axpy_opencl.c

@@ -0,0 +1,73 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2012 inria
+ *
+ * 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_opencl.h>
+
+extern struct starpu_opencl_program opencl_program;
+
+void axpy_opencl(void *buffers[], void *_args)
+{
+	float *alpha = _args;
+	int id, devid;
+        cl_int err;
+	cl_kernel kernel;
+	cl_command_queue queue;
+	cl_event event;
+
+	unsigned n = STARPU_VECTOR_GET_NX(buffers[0]);
+	cl_mem x = (cl_mem) STARPU_VECTOR_GET_DEV_HANDLE(buffers[0]);
+	cl_mem y = (cl_mem) STARPU_VECTOR_GET_DEV_HANDLE(buffers[1]);
+
+	id = starpu_worker_get_id();
+	devid = starpu_worker_get_devid(id);
+
+	err = starpu_opencl_load_kernel(&kernel, &queue, &opencl_program, "_axpy_opencl", devid);
+	if (err != CL_SUCCESS)
+		STARPU_OPENCL_REPORT_ERROR(err);
+
+	err = clSetKernelArg(kernel, 0, sizeof(x), &x);
+	err|= clSetKernelArg(kernel, 1, sizeof(y), &y);
+	err|= clSetKernelArg(kernel, 2, sizeof(n), &n);
+	err|= clSetKernelArg(kernel, 3, sizeof(*alpha), alpha);
+	if (err)
+		STARPU_OPENCL_REPORT_ERROR(err);
+
+	{
+		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;
+
+		err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, &local, 0, NULL, &event);
+		if (err != CL_SUCCESS)
+			STARPU_OPENCL_REPORT_ERROR(err);
+	}
+
+	clFinish(queue);
+	starpu_opencl_collect_stats(event);
+	clReleaseEvent(event);
+
+	starpu_opencl_release_kernel(kernel);
+}

examples/axpy/axpy_opencl_kernel.cl

@@ -0,0 +1,27 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2012 inria
+ *
+ * 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 "axpy.h"
+
+__kernel void _axpy_opencl(__global TYPE *x,
+			   __global TYPE *y,
+			   unsigned nx,
+			   TYPE alpha)
+{
+        const int i = get_global_id(0);
+        if (i < nx)
+                y[i] = alpha * x[i] + y[i];
+}