gcc: Add an OpenCL task implementation for `vector_scal'.

* gcc-plugin/examples/vector_scal/vector_scal.c (cl_programs): New
  variable.
  (vector_scal_opencl): New task implementation.
  (main)[STARPU_USE_OPENCL]: Load `vector_scal_opencl_kernel.cl'.

* gcc-plugin/examples/vector_scal/vector_scal_opencl_kernel.cl: New file.

* gcc-plugin/examples/Makefile.am (EXTRA_DIST): New variable.

.gitignore

@@ -188,3 +188,4 @@ starpu.log
 /gcc-plugin/src/c-expr.c
 /gcc-plugin/tests/heap-allocated
 /gcc-plugin/tests/output-pointer
+/gcc-plugin/examples/vector_scal/vector_scal

gcc-plugin/examples/Makefile.am

@@ -40,6 +40,8 @@ noinst_HEADERS =				\
   cholesky/cholesky.h				\
   cholesky/cholesky_kernels.h
 
+EXTRA_DIST = vector_scal/vector_scal_opencl_kernel.cl
+
 if !NO_BLAS_LIB
 cholesky_cholesky_SOURCES	=		\
 	cholesky/cholesky.c		\

gcc-plugin/examples/vector_scal/vector_scal.c

@@ -26,12 +26,12 @@
 
    Features to test in a near future :
 	- CUDA
-	- OpenCL
 	- Filters
  */
 #include <stdio.h>
 #include <stdlib.h>
 
+
 /* Declare and define the standard CPU implementation.  */
 
 static void vector_scal (size_t size, float vector[size], float factor)
@@ -48,6 +48,7 @@ vector_scal_cpu (size_t size, float vector[size], float factor)
     vector[i] *= factor;
 }
 
+
 #ifdef __SSE__
 /* The SSE-capable CPU implementation.  */
 
@@ -79,6 +80,59 @@ vector_scal_sse (size_t size, float vector[size], float factor)
 }
 #endif /* __SSE__ */
 
+
+/* Declaration and definition of the OpenCL implementation.  */
+
+#ifdef STARPU_USE_OPENCL
+
+/* The OpenCL programs, loaded from `main'.  */
+static struct starpu_opencl_program cl_programs;
+
+static void vector_scal_opencl (size_t size, float vector[size], float factor)
+  __attribute__ ((task_implementation ("opencl", vector_scal)));
+
+static void
+vector_scal_opencl (size_t size, float vector[size], float factor)
+{
+  int id, devid, err;
+  cl_kernel kernel;
+  cl_command_queue queue;
+  cl_event event;
+
+  cl_mem val = (cl_mem) vector;
+
+  id = starpu_worker_get_id ();
+  devid = starpu_worker_get_devid (id);
+
+  /* Prepare to invoke the kernel.  In the future, this will be largely
+     automated.  */
+  err = starpu_opencl_load_kernel (&kernel, &queue, &programs,
+				   "vector_mult_opencl", devid);
+  if (err != CL_SUCCESS)
+    STARPU_OPENCL_REPORT_ERROR (err);
+
+  err = clSetKernelArg (kernel, 0, sizeof (val), &val);
+  err |= clSetKernelArg (kernel, 1, sizeof (size), &size);
+  err |= clSetKernelArg (kernel, 2, sizeof (factor), &factor);
+  if (err)
+    STARPU_OPENCL_REPORT_ERROR (err);
+
+  size_t global = 1, local = 1;
+  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);
+}
+
+#endif
+
+
 #define EPSILON 1e-3
 static int
 check (size_t size, float vector[size], float factor)
@@ -97,6 +151,11 @@ main (void)
 {
 #pragma starpu initialize
 
+#ifdef STARPU_USE_OPENCL
+  starpu_opencl_load_opencl_from_file ("vector_scal_opencl_kernel.cl",
+				       &cl_programs, NULL);
+#endif
+
 #define NX     0x100000
 #define FACTOR 3.14
 

gcc-plugin/examples/vector_scal/vector_scal_opencl_kernel.cl

@@ -0,0 +1,38 @@
+/* StarPU --- Runtime system for heterogeneous multicore architectures.
+ *
+ * Copyright (C) 2010, 2011  Centre National de la Recherche Scientifique
+ * Copyright (C) 2010, 2011  Université de Bordeaux 1
+ *
+ * Redistribution  and  use  in  source and binary forms, with or without
+ * modification,  are  permitted  provided  that the following conditions
+ * are met:
+ *
+ * * Redistributions  of  source  code  must  retain  the above copyright
+ *   notice,  this  list  of  conditions  and  the  following  disclaimer.
+ * * Redistributions  in  binary  form must reproduce the above copyright
+ *   notice,  this list of conditions and the following disclaimer in the
+ *   documentation  and/or other materials provided with the distribution.
+ * * The name of the author may not be used to endorse or promote products
+ *   derived from this software without specific prior written permission.
+ *
+ * THIS  SOFTWARE  IS  PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
+ * ``AS IS''  AND  ANY  EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
+ * LIMITED  TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR
+ * A  PARTICULAR  PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT
+ * HOLDERS OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL
+ * SPECIAL,  EXEMPLARY,  OR  CONSEQUENTIAL  DAMAGES  (INCLUDING,  BUT NOT
+ * LIMITED  TO,  PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE
+ * DATA,  OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY
+ * THEORY  OF  LIABILITY,  WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING  NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ * OF  THIS  SOFTWARE,  EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+
+__kernel void vector_mult_opencl(__global float* val, int nx, float factor)
+{
+        const int i = get_global_id(0);
+        if (i < nx) {
+                val[i] *= factor;
+        }
+}
+