gcc-plugin: added groupsize support to OpenCL codelet generation.

* gcc-plugin/src/starpu.c (define_opencl_task_implementation):
  added new function argument and check its type.
  (handle_pragma_opencl): added support for the groupsize pragma argument.
* gcc-plugin/tests/mocks.h (clEnqueueNDRangeKernel):
  Completed mock function.
* gcc-plugin/tests/opencl-errors.c: added new tests.
* gcc-plugin/tests/opencl.c: completed test.
  (main): updated check to the `clEnqueueNDRangeKernel' function call.

gcc-plugin/src/starpu.c

@@ -639,7 +639,7 @@ handle_pragma_initialize (struct cpp_reader *reader)
 			      error_var, integer_zero_node),
 		      build_error_statements (loc, error_var,
 					      build_starpu_error_string,
-					      "failed to initialize StarPU"), 
+					      "failed to initialize StarPU"),
 		      NULL_TREE);
 
   tree stmts = NULL_TREE;
@@ -1278,11 +1278,13 @@ build_opencl_set_kernel_arg_calls (location_t loc, tree task_impl,
 
 static void
 define_opencl_task_implementation (location_t loc, tree task_impl,
-				   const char *file, const_tree kernel)
+				   const char *file, const_tree kernel,
+				   const_tree groupsize)
 {
   gcc_assert (task_implementation_p (task_impl)
 	      && task_implementation_where (task_impl) == STARPU_OPENCL);
   gcc_assert (TREE_CODE (kernel) == STRING_CST);
+  gcc_assert (INTEGRAL_TYPE_P (TREE_TYPE (groupsize)));
 
   local_define (tree, local_var, (tree type))
   {
@@ -1452,8 +1454,7 @@ define_opencl_task_implementation (location_t loc, tree task_impl,
 
       /* TODO: Support user-provided values.  */
       append_to_statement_list (build2 (INIT_EXPR, TREE_TYPE (group_size_var),
-					group_size_var,
-					build_int_cst (integer_type_node, 1)),
+					group_size_var, (tree)groupsize),
 				&stmts);
       append_to_statement_list (build2 (INIT_EXPR, TREE_TYPE (ngroups_var),
 					ngroups_var,
@@ -1532,8 +1533,8 @@ handle_pragma_opencl (struct cpp_reader *reader)
   if (args == NULL_TREE)
     return;
 
-  /* TODO: Add "group size" and "number of groups" arguments.  */
-  if (list_length (args) < 3)
+  /* TODO: Add "number of groups" arguments.  */
+  if (list_length (args) < 4)
     {
       error_at (loc, "wrong number of arguments for %<starpu opencl%> pragma");
       return;
@@ -1552,13 +1553,21 @@ handle_pragma_opencl (struct cpp_reader *reader)
   	      if (TREE_CODE (TREE_VALUE (args)) == STRING_CST)
   		{
   		  tree kernel = TREE_VALUE (args);
+		  args = TREE_CHAIN (args);
 
-  		  if (TREE_CHAIN (args) == NULL_TREE)
-		    define_opencl_task_implementation (loc, task_impl,
-						       TREE_STRING_POINTER (file),
-						       kernel);
-  		  else
-  		    error_at (loc, "junk after %<starpu opencl%> pragma");
+		  if (TREE_TYPE (TREE_VALUE (args)) != NULL_TREE &&
+		      INTEGRAL_TYPE_P (TREE_TYPE (TREE_VALUE (args))))
+		    {
+		      tree groupsize = TREE_VALUE (args);
+		      if (TREE_CHAIN (args) == NULL_TREE)
+			define_opencl_task_implementation (loc, task_impl,
+							   TREE_STRING_POINTER (file),
+							   kernel, groupsize);
+		      else
+			error_at (loc, "junk after %<starpu opencl%> pragma");
+		    }
+		  else
+		    error_at (loc, "%<groupsize%> argument must be an integral type");
   		}
   	      else
   		error_at (loc, "%<kernel%> argument must be a string constant");

gcc-plugin/tests/mocks.h

@@ -469,6 +469,15 @@ struct load_opencl_arguments
 /* Expected arguments.  */
 static struct load_opencl_arguments expected_load_opencl_arguments;
 
+struct cl_enqueue_kernel_arguments
+{
+  size_t * global_work_size;
+};
+
+/* Variable describing the expected `clEnqueueNDRangeKernel' arguments. */
+static struct cl_enqueue_kernel_arguments expected_cl_enqueue_kernel_arguments;
+
+
 int
 starpu_opencl_load_opencl_from_string (const char *source,
 				       struct starpu_opencl_program *program,
@@ -551,7 +560,9 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue,
                        const cl_event * event_wait_list,
                        cl_event *       event)
 {
-  assert (*global_work_size == 1 && *local_work_size == 1);
+  assert (*local_work_size == 1);
+  assert (*global_work_size == *expected_cl_enqueue_kernel_arguments.global_work_size);
+
   opencl_enqueue_calls++;
   return 0;
 }

gcc-plugin/tests/opencl-errors.c

@@ -35,19 +35,20 @@ my_task_cpu (int x, float a[x])
 }
 
 
-#pragma starpu opencl my_task "test.cl" "kern" /* (error "not a.* task impl") */
+#pragma starpu opencl my_task "test.cl" "kern" 1 /* (error "not a.* task impl") */
 #pragma starpu opencl my_task_cpu  /* (error "not a.* task impl") */	\
-                      "test.cl" "kern"
-#pragma starpu opencl my_task_opencl "/dev/null" "kern" /* (error "empty") */
-#pragma starpu opencl my_task_opencl "/does-not-exist/" "kern" /* (error "failed to access") */
+                      "test.cl" "kern" 1
+#pragma starpu opencl my_task_opencl "/dev/null" "kern" 1 /* (error "empty") */
+#pragma starpu opencl my_task_opencl "/does-not-exist/" "kern" 1 /* (error "failed to access") */
 
 #pragma starpu opencl my_task_opencl	  /* (error "wrong number of arg") */
-#pragma starpu opencl my_task_opencl 123 "kern" /* (error "string constant") */
-#pragma starpu opencl my_task_opencl "test.cl" 123 /* (error "string constant") */
-#pragma starpu opencl my_task_opencl "test.cl" "kern" "foo" /* (error "junk after") */
+#pragma starpu opencl my_task_opencl 123 "kern" 1 /* (error "string constant") */
+#pragma starpu opencl my_task_opencl "test.cl" 123 1 /* (error "string constant") */
+#pragma starpu opencl my_task_opencl "test.cl" "kern" "a" /* (error "integral type") */
+#pragma starpu opencl my_task_opencl "test.cl" "kern" 1 "foo" /* (error "junk after") */
 
 void
 foo (void)
 {
-#pragma starpu opencl my_task_opencl "test.cl" "kern" /* (error "top-level") */
+#pragma starpu opencl my_task_opencl "test.cl" "kern" 1 /* (error "top-level") */
 }

gcc-plugin/tests/opencl.c

@@ -29,7 +29,7 @@ static void my_task (int x, float a[x])
 static void my_task_opencl (int x, float a[x])
   __attribute__ ((task_implementation ("opencl", my_task)));
 
-#pragma starpu opencl my_task_opencl "test.cl" "kern"
+#pragma starpu opencl my_task_opencl "test.cl" "kern" 8
 
 int
 main ()
@@ -55,6 +55,7 @@ main ()
 
   expected_insert_task_arguments = expected;
   expected_insert_task_targets = STARPU_OPENCL;
+  size_t y = 8; expected_cl_enqueue_kernel_arguments.global_work_size = &y;
 
   my_task (123, a);
   my_task (123, a);