gcc: Error out when a task with an OpenCL impl. uses `size_t' arguments.

* gcc-plugin/src/starpu.c (validate_opencl_argument_type): New procedure.

* gcc-plugin/tests/opencl-size_t.c: New file.
* gcc-plugin/tests/Makefile.am (gcc_tests): Add `opencl-size_t.c'.

* gcc-plugin/tests/output-pointer.c: Use `int' instead of `size_t' for
  task arguments.

* gcc-plugin/examples/vector_scal/vector_scal.c (vector_scal,
  vector_scal_cpu, vector_scal_sse, vector_scal_opencl,
  vector_scal_cuda): Use `unsigned int' instead of `size_t'.
* gcc-plugin/examples/vector_scal/vector_scal_opencl_kernel.cl:
  Likewise.

* doc/chapters/basic-examples.texi (Hello World using the C Extension):
  Change examples to use `unsigned' instead of `size_t'.
* doc/chapters/c-extensions.texi: Likewise.

doc/chapters/basic-examples.texi

@@ -68,7 +68,8 @@ has a single implementation for CPU:
 static void my_task (int x) __attribute__ ((task));
 
 /* Declaration of the CPU implementation of `my_task'.  */
-static void my_task_cpu (int x) __attribute__ ((task_implementation ("cpu", my_task)));
+static void my_task_cpu (int x)
+   __attribute__ ((task_implementation ("cpu", my_task)));
 
 /* Definition of said CPU implementation.  */
 static void my_task_cpu (int x)
@@ -327,20 +328,20 @@ has to be defined:
 @smallexample
 /* Declare the `vector_scal' task.  */
 
-static void vector_scal (size_t size, float vector[size],
+static void vector_scal (unsigned size, float vector[size],
                          float factor)
   __attribute__ ((task));
 
 /* Declare and define the standard CPU implementation.  */
 
-static void vector_scal_cpu (size_t size, float vector[size],
+static void vector_scal_cpu (unsigned size, float vector[size],
                              float factor)
   __attribute__ ((task_implementation ("cpu", vector_scal)));
 
 static void
-vector_scal_cpu (size_t size, float vector[size], float factor)
+vector_scal_cpu (unsigned size, float vector[size], float factor)
 @{
-  size_t i;
+  unsigned i;
   for (i = 0; i < size; i++)
     vector[i] *= factor;
 @}
@@ -447,13 +448,12 @@ in our C file like this:
 /* The OpenCL programs, loaded from `main' (see below).  */
 static struct starpu_opencl_program cl_programs;
 
-static void vector_scal_opencl (size_t size, float vector[size],
+static void vector_scal_opencl (unsigned size, float vector[size],
                                 float factor)
   __attribute__ ((task_implementation ("opencl", vector_scal)));
 
-@c TODO This example will not work : size cannot be a size_t in clSetKernelArg, and global should not be 1. Do we want to document the ugly hach we use, though ?
 static void
-vector_scal_opencl (size_t size, float vector[size], float factor)
+vector_scal_opencl (unsigned size, float vector[size], float factor)
 @{
   int id, devid, err;
   cl_kernel kernel;
@@ -523,7 +523,7 @@ declaration for the task implementation:
 
 @cartouche
 @smallexample
-extern void vector_scal_cuda (size_t size, float vector[size],
+extern void vector_scal_cuda (unsigned size, float vector[size],
                               float factor)
   __attribute__ ((task_implementation ("cuda", vector_scal)));
 @end smallexample

doc/chapters/c-extensions.texi

@@ -120,20 +120,20 @@ Here is an example:
 
 static void matmul (const float *A, const float *B,
                     __output float *C,
-                    size_t nx, size_t ny, size_t nz)
+                    unsigned nx, unsigned ny, unsigned nz)
   __attribute__ ((task));
 
 static void matmul_cpu (const float *A, const float *B,
                         __output float *C,
-                        size_t nx, size_t ny, size_t nz)
+                        unsigned nx, unsigned ny, unsigned nz)
   __attribute__ ((task_implementation ("cpu", matmul)));
 
 
 static void
 matmul_cpu (const float *A, const float *B, __output float *C,
-            size_t nx, size_t ny, size_t nz)
+            unsigned nx, unsigned ny, unsigned nz)
 @{
-  size_t i, j, k;
+  unsigned i, j, k;
 
   for (j = 0; j < ny; j++)
     for (i = 0; i < nx; i++)
@@ -156,11 +156,11 @@ CUDA and OpenCL implementations can be declared in a similar way:
 @cartouche
 @smallexample
 static void matmul_cuda (const float *A, const float *B, float *C,
-                         size_t nx, size_t ny, size_t nz)
+                         unsigned nx, unsigned ny, unsigned nz)
   __attribute__ ((task_implementation ("cuda", matmul)));
 
 static void matmul_opencl (const float *A, const float *B, float *C,
-                           size_t nx, size_t ny, size_t nz)
+                           unsigned nx, unsigned ny, unsigned nz)
   __attribute__ ((task_implementation ("opencl", matmul)));
 @end smallexample
 @end cartouche
@@ -175,7 +175,7 @@ OpenCL under the hood, such as CUBLAS functions:
 @smallexample
 static void
 matmul_cuda (const float *A, const float *B, float *C,
-             size_t nx, size_t ny, size_t nz)
+             unsigned nx, unsigned ny, unsigned nz)
 @{
   cublasSgemm ('n', 'n', nx, ny, nz,
                1.0f, A, 0, B, 0,
@@ -356,12 +356,12 @@ in a way that allows it to be compiled without the GCC plug-in:
 #include <stdlib.h>
 
 static void matmul (const float *A, const float *B, float *C,
-                    size_t nx, size_t ny, size_t nz) __task;
+                    unsigned nx, unsigned ny, unsigned nz) __task;
 
 #ifdef STARPU_GCC_PLUGIN
 
 static void matmul_cpu (const float *A, const float *B, float *C,
-                        size_t nx, size_t ny, size_t nz)
+                        unsigned nx, unsigned ny, unsigned nz)
   __attribute__ ((task_implementation ("cpu", matmul)));
 
 #endif
@@ -369,7 +369,7 @@ static void matmul_cpu (const float *A, const float *B, float *C,
 
 static void
 CPU_TASK_IMPL (matmul) (const float *A, const float *B, float *C,
-                        size_t nx, size_t ny, size_t nz)
+                        unsigned nx, unsigned ny, unsigned nz)
 @{
   /* Code of the CPU kernel here...  */
 @}

gcc-plugin/examples/vector_scal/vector_scal.c

@@ -29,16 +29,16 @@
 
 /* Declare and define the standard CPU implementation.  */
 
-static void vector_scal (size_t size, float vector[size], float factor)
+static void vector_scal (unsigned int size, float vector[size], float factor)
   __attribute__ ((task));
 
-static void vector_scal_cpu (size_t size, float vector[size], float factor)
+static void vector_scal_cpu (unsigned int size, float vector[size], float factor)
   __attribute__ ((task_implementation ("cpu", vector_scal)));
 
 static void
-vector_scal_cpu (size_t size, float vector[size], float factor)
+vector_scal_cpu (unsigned int size, float vector[size], float factor)
 {
-  size_t i;
+  unsigned int i;
   for (i = 0; i < size; i++)
     vector[i] *= factor;
 }
@@ -49,11 +49,11 @@ vector_scal_cpu (size_t size, float vector[size], float factor)
 
 #include <xmmintrin.h>
 
-static void vector_scal_sse (size_t size, float vector[size], float factor)
+static void vector_scal_sse (unsigned int size, float vector[size], float factor)
   __attribute__ ((task_implementation ("cpu", vector_scal)));
 
 static void
-vector_scal_sse (size_t size, float vector[size], float factor)
+vector_scal_sse (unsigned int size, float vector[size], float factor)
 {
   unsigned int n_iterations = size / 4;
 
@@ -85,11 +85,11 @@ vector_scal_sse (size_t size, float vector[size], float factor)
 /* 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)
+static void vector_scal_opencl (unsigned int 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)
+vector_scal_opencl (unsigned int size, float vector[size], float factor)
 {
   int id, devid, err;
   cl_kernel kernel;
@@ -108,15 +108,13 @@ vector_scal_opencl (size_t size, float vector[size], float factor)
   if (err != CL_SUCCESS)
     STARPU_OPENCL_REPORT_ERROR (err);
 
-  /* XXX : clSetKernelArg will not work with a size_t ... */
-  int _size = size;
   err = clSetKernelArg (kernel, 0, sizeof (val), &val);
-  err |= clSetKernelArg (kernel, 1, sizeof (_size), &_size);
+  err |= clSetKernelArg (kernel, 1, sizeof (size), &size);
   err |= clSetKernelArg (kernel, 2, sizeof (factor), &factor);
   if (err)
     STARPU_OPENCL_REPORT_ERROR (err);
 
-  size_t global = _size, local = 1;
+  size_t global = 1, local = 1;
   err = clEnqueueNDRangeKernel (queue, kernel, 1, NULL, &global, &local, 0,
 				NULL, &event);
   if (err != CL_SUCCESS)
@@ -137,7 +135,7 @@ vector_scal_opencl (size_t size, float vector[size], float factor)
 /* Declaration of the CUDA implementation.  The definition itself is in the
    `.cu' file itself.  */
 
-extern void vector_scal_cuda (size_t size, float vector[size], float factor)
+extern void vector_scal_cuda (unsigned int size, float vector[size], float factor)
   __attribute__ ((task_implementation ("cuda", vector_scal)));
 
 #endif

gcc-plugin/examples/vector_scal/vector_scal_opencl_kernel.cl

@@ -28,7 +28,7 @@
  * OF  THIS  SOFTWARE,  EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
  */
 
-__kernel void vector_mult_opencl(__global float* val, int nx, float factor)
+__kernel void vector_mult_opencl(__global float* val, unsigned int nx, float factor)
 {
         const int i = get_global_id(0);
         if (i < nx) {

gcc-plugin/src/starpu.c

@@ -995,6 +995,24 @@ handle_task_attribute (tree *node, tree name, tree args,
   return NULL_TREE;
 }
 
+/* Diagnose use of C types that are either nonexistent or different in
+   OpenCL.  */
+
+static void
+validate_opencl_argument_type (location_t loc, const_tree type)
+{
+  if (INTEGRAL_TYPE_P (type))
+    {
+      tree decl = TYPE_NAME (type);
+
+      if (DECL_P (decl)
+	  && DECL_IN_SYSTEM_HEADER (decl)
+	  && DECL_NAME (decl) == get_identifier ("size_t"))
+	/* Check for the use of a literal `size_t'.  */
+	error_at (loc, "%<size_t%> is not a valid OpenCL type");
+    }
+}
+
 /* Handle the `task_implementation (WHERE, TASK)' attribute.  WHERE is a
    string constant ("cpu", "cuda", etc.), and TASK is the identifier of a
    function declared with the `task' attribute.  */
@@ -1059,6 +1077,15 @@ handle_task_implementation_attribute (tree *node, tree name, tree args,
 	warning_at (loc, 0,
 		    "unsupported target %E; task implementation won't be used",
 		    where);
+      else if (task_implementation_target_to_int (where) == STARPU_OPENCL)
+	{
+	  void validate (tree t)
+	  {
+	    validate_opencl_argument_type (loc, t);
+	  }
+
+	  for_each (validate, TYPE_ARG_TYPES (TREE_TYPE (fn)));
+	}
 
       /* Keep the attribute.  */
       *no_add_attrs = false;

gcc-plugin/tests/Makefile.am

@@ -36,6 +36,7 @@ gcc_tests =					\
   heap-allocated-errors.c			\
   verbose.c					\
   debug-tree.c					\
+  opencl-size_t.c				\
   shutdown-errors.c
 
 dist_noinst_HEADERS = mocks.h

gcc-plugin/tests/opencl-size_t.c

@@ -0,0 +1,39 @@
+/* GCC-StarPU
+   Copyright (C) 2012 Institut National de Recherche en Informatique et Automatique
+
+   GCC-StarPU is free software: you can redistribute it and/or modify
+   it under the terms of the GNU General Public License as published by
+   the Free Software Foundation, either version 3 of the License, or
+   (at your option) any later version.
+
+   GCC-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 General Public License for more details.
+
+   You should have received a copy of the GNU General Public License
+   along with GCC-StarPU.  If not, see <http://www.gnu.org/licenses/>.  */
+
+/* Make sure use of `size_t' as a task argument type is flagged.  */
+
+#undef NDEBUG
+
+#include <mocks.h>
+#include <unistd.h>
+
+static void my_task (size_t size, int x[size]) __attribute__ ((task));
+
+static void my_task_cpu (size_t size, int x[size])
+  __attribute__ ((task_implementation ("cpu", my_task)));
+static void my_task_opencl (size_t size, int x[size]) /* (error "not a valid OpenCL type") */
+  __attribute__ ((task_implementation ("opencl", my_task)));
+
+static void
+my_task_cpu (size_t size, int x[size])
+{
+}
+
+static void
+my_task_opencl (size_t size, int x[size])
+{
+}

gcc-plugin/tests/output-pointer.c

@@ -22,46 +22,46 @@
 
 /* The tasks under test.  */
 
-static void my_pointer_task (size_t size, __output int *x)
+static void my_pointer_task (int size, __output int *x)
   __attribute__ ((task));
 
-static void my_pointer_task_cpu (size_t size, __output int *x)
+static void my_pointer_task_cpu (int size, __output int *x)
   __attribute__ ((task_implementation ("cpu", my_pointer_task)));
-static void my_pointer_task_opencl (size_t size, __output int *x)
+static void my_pointer_task_opencl (int size, __output int *x)
   __attribute__ ((task_implementation ("opencl", my_pointer_task)));
 
 static void
-my_pointer_task_cpu (size_t size, __output int *x)
+my_pointer_task_cpu (int size, __output int *x)
 {
-  printf ("%s: x = %p, size = %zi\n", __func__, x, size);
+  printf ("%s: x = %p, size = %i\n", __func__, x, size);
 }
 
 static void
-my_pointer_task_opencl (size_t size, int *x)
+my_pointer_task_opencl (int size, int *x)
 {
-  printf ("%s: x = %p, size = %zi\n", __func__, x, size);
+  printf ("%s: x = %p, size = %i\n", __func__, x, size);
 }
 
 
 
-static void my_array_task (size_t size, __output int x[size])
+static void my_array_task (int size, __output int x[size])
   __attribute__ ((task));
 
-static void my_array_task_cpu (size_t size, __output int x[size])
+static void my_array_task_cpu (int size, __output int x[size])
   __attribute__ ((task_implementation ("cpu", my_array_task)));
-static void my_array_task_opencl (size_t size, __output int x[size])
+static void my_array_task_opencl (int size, __output int x[size])
   __attribute__ ((task_implementation ("opencl", my_array_task)));
 
 static void
-my_array_task_cpu (size_t size, __output int x[size])
+my_array_task_cpu (int size, __output int x[size])
 {
-  printf ("%s: x = %p, size = %zi\n", __func__, x, size);
+  printf ("%s: x = %p, size = %i\n", __func__, x, size);
 }
 
 static void
-my_array_task_opencl (size_t size, __output int x[size])
+my_array_task_opencl (int size, __output int x[size])
 {
-  printf ("%s: x = %p, size = %zi\n", __func__, x, size);
+  printf ("%s: x = %p, size = %i\n", __func__, x, size);
 }
 
 
@@ -71,7 +71,7 @@ main (int argc, char *argv[])
 {
 #pragma starpu initialize
 
-  size_t size = 42;
+  int size = 42;
   int x[size];
 
   /* Register X (don't use the pragma, to avoid mixing concerns in this