example: add OpenCL codelet for vector scaling example

examples/Makefile.am

@@ -30,6 +30,7 @@ check_PROGRAMS =
 BUILT_SOURCES =
 
 EXTRA_DIST = 					\
+	basic_examples/vector_scal_opencl_codelet.cl \
 	cuda/incrementer_cuda.cu		\
 	spmv/spmv_cuda.cu			\
 	gordon/null_kernel_gordon.c		\
@@ -146,6 +147,14 @@ basic_examples_vector_scal_SOURCES +=		\
 	basic_examples/vector_scal_cuda.cu
 endif
 
+if STARPU_USE_OPENCL
+basic_examples_vector_scal_SOURCES +=		\
+	basic_examples/vector_scal_opencl.c
+endif
+
+nobase_STARPU_OPENCL_DATA_DATA = \
+	basic_examples/vector_scal_opencl_codelet.cl
+
 examplebin_PROGRAMS +=				\
 	basic_examples/mult
 
@@ -472,9 +481,11 @@ endif
 if STARPU_USE_OPENCL
 variable_variable_SOURCES +=	\
 	variable/variable_kernels_opencl.c
+endif
+
 nobase_STARPU_OPENCL_DATA_DATA += \
 	variable/variable_kernels_opencl_codelet.cl
-endif
+
 
 ######################
 # matVecMult example #

examples/basic_examples/vector_scal.c

@@ -62,12 +62,16 @@ static void scal_cpu_func(void *buffers[], void *cl_arg)
 }
 
 extern void scal_cuda_func(void *buffers[], void *_args);
+extern void scal_opencl_func(void *buffers[], void *_args);
 
 static starpu_codelet cl = {
 	.where = STARPU_CPU
 #ifdef STARPU_USE_CUDA
 		| STARPU_CUDA
 #endif
+#ifdef STARPU_USE_OPENCL
+		| STARPU_OPENCL
+#endif
 		,
 	/* CPU implementation of the codelet */
 	.cpu_func = scal_cpu_func,
@@ -75,6 +79,10 @@ static starpu_codelet cl = {
 	/* CUDA implementation of the codelet */
 	.cuda_func = scal_cuda_func,
 #endif
+#ifdef STARPU_USE_OPENCL
+	/* OpenCL implementation of the codelet */
+	.opencl_func = scal_opencl_func,
+#endif
 	.nbuffers = 1
 };
 
@@ -92,6 +100,10 @@ int main(int argc, char **argv)
 	/* Initialize StarPU with default configuration */
 	starpu_init(NULL);
 
+#ifdef STARPU_USE_OPENCL
+        _starpu_opencl_compile_source_to_opencl("examples/basic_examples/vector_scal_opencl_codelet.cl");
+#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.

examples/basic_examples/vector_scal_opencl.c

@@ -0,0 +1,62 @@
+/*
+ * 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.
+ */
+
+/*
+ * This example complements vector_scale.c: here we implement a OpenCL version.
+ */
+
+#include <starpu.h>
+#include <starpu_opencl.h>
+
+void scal_opencl_func(void *buffers[], void *_args)
+{
+        float *factor = (float *)_args;
+        struct starpu_vector_interface_s *vector = (struct starpu_vector_interface_s *) buffers[0];
+	int id, devid, err;
+	cl_kernel kernel;
+	cl_command_queue queue;
+
+        /* length of the vector */
+        unsigned n = STARPU_GET_VECTOR_NX(vector);
+        /* local copy of the vector pointer */
+        float *val = (float *)STARPU_GET_VECTOR_PTR(vector);
+
+        id = starpu_worker_get_id();
+        devid = starpu_worker_get_devid(id);
+
+	err = starpu_opencl_load_kernel(&kernel, &queue,
+                                        "examples/basic_examples/vector_scal_opencl_codelet.cl", "vectorScal", devid);
+	if (err != CL_SUCCESS) STARPU_OPENCL_REPORT_ERROR(err);
+
+	err = 0;
+	err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &val);
+	err = clSetKernelArg(kernel, 1, sizeof(int), &n);
+        err |= clSetKernelArg(kernel, 2, sizeof(float), (void*)factor);
+
+	if (err) STARPU_OPENCL_REPORT_ERROR(err);
+
+	{
+		size_t global=1;
+		size_t local=1;
+		err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, &local, 0, NULL, NULL);
+		if (err != CL_SUCCESS) STARPU_OPENCL_REPORT_ERROR(err);
+	}
+
+	clFinish(queue);
+
+	starpu_opencl_release(kernel);
+
+}

examples/basic_examples/vector_scal_opencl_codelet.cl

@@ -0,0 +1,22 @@
+/*
+ * Copyright 1993-2009 NVIDIA Corporation.  All rights reserved.
+ *
+ * NVIDIA Corporation and its licensors retain all intellectual property and
+ * proprietary rights in and to this software and related documentation.
+ * Any use, reproduction, disclosure, or distribution of this software
+ * and related documentation without an express license agreement from
+ * NVIDIA Corporation is strictly prohibited.
+ *
+ * Please refer to the applicable NVIDIA end user license agreement (EULA)
+ * associated with this source code for terms and conditions that govern
+ * your use of this NVIDIA software.
+ *
+ */
+
+__kernel void vectorScal(__global float* val, int nx, float factor)
+{
+        const int i = get_global_id(0);
+        if (i < nx) {
+                val[i] *= factor;
+        }
+}