merge trunk@6588:6595

configure.ac

@@ -1529,7 +1529,10 @@ AM_CONDITIONAL(BUILD_STARPUFFT, [test x$fft_support = xyes])
 # hwloc                                  #
 ##########################################
 
-AC_ARG_WITH([hwloc], [AS_HELP_STRING([--without-hwloc], [Disable hwloc (enabled by default)])])
+AC_ARG_WITH([hwloc],
+	[AS_HELP_STRING([--without-hwloc],
+	[Disable hwloc (enabled by default)])],
+	[hwloc_dir="$withval"])
 SAVED_LDFLAGS="${LDFLAGS}"
 SAVED_CPPFLAGS="${CPPFLAGS}"
 AS_IF([test "x$with_hwloc" != "xno"], [
@@ -1539,10 +1542,9 @@ AS_IF([test "x$with_hwloc" != "xno"], [
 		have_valid_hwloc=no
 		have_pkgconfig_hwloc=no])
 	AS_IF([test "$have_valid_hwloc" = "no"], [
-		hwloc_dir="$withval"
-		CPPFLAGS="${SAVED_CPPFLAGS} -I$hwloc_dir/include"
+		if test "$hwloc_dir" != "" ; then CPPFLAGS="${SAVED_CPPFLAGS} -I$hwloc_dir/include" ; fi
 		AC_CHECK_HEADER([hwloc.h],[have_valid_hwloc=yes],[have_valid_hwloc=no])
-		LDFLAGS="${SAVED_LDFLAGS} -L$hwloc_dir/lib"
+		if test "$hwloc_dir" != "" ; then LDFLAGS="${SAVED_LDFLAGS} -L$hwloc_dir/lib" ; fi
 		AC_HAVE_LIBRARY([hwloc],[have_valid_hwloc=yes],[have_valid_hwloc=no])
 		])
     ],

doc/chapters/basic-api.texi

@@ -2303,6 +2303,11 @@ collect statistics about the kernel execution (used cycles, consumed power).
 @node OpenCL utilities
 @subsection OpenCL utilities
 
+@deftypefun {const char *}starpu_opencl_error_string (cl_int @var{status})
+Return the error message in English corresponding to @var{status}, an
+OpenCL error code.
+@end deftypefun
+
 @deftypefun void starpu_opencl_display_error ({const char *}@var{func}, {const char *}@var{file}, int @var{line}, {const char *}@var{msg}, cl_int @var{status})
 Given a valid error @var{status}, prints the corresponding error message on
 stdout, along with the given function name @var{func}, the given filename

gcc-plugin/src/starpu.c

@@ -412,13 +412,18 @@ build_hello_world (void)
 }
 
 /* Given ERROR_VAR, an integer variable holding a StarPU error code, return
-   statements that print out an error message and abort.  */
+   statements that print out the error message returned by
+   BUILD_ERROR_MESSAGE (ERROR_VAR) and abort.  */
 
-static tree build_error_statements (location_t, tree, const char *, ...)
-  __attribute__ ((format (printf, 3, 4)));
+static tree build_error_statements (location_t, tree,
+				    function_parm (tree, f, (tree)),
+				    const char *, ...)
+  __attribute__ ((format (printf, 4, 5)));
 
 static tree
-build_error_statements (location_t loc, tree error_var, const char *fmt, ...)
+build_error_statements (location_t loc, tree error_var,
+			function_parm (tree, build_error_message, (tree)),
+			const char *fmt, ...)
 {
   expanded_location xloc = expand_location (loc);
 
@@ -436,23 +441,17 @@ build_error_statements (location_t loc, tree error_var, const char *fmt, ...)
   if (error_var != NULL_TREE)
     {
       /* ERROR_VAR is an error code.  */
-
-      static tree strerror_fn;
-      LOOKUP_STARPU_FUNCTION (strerror_fn, "strerror");
-
       gcc_assert (TREE_CODE (error_var) == VAR_DECL
 		  && TREE_TYPE (error_var) == integer_type_node);
 
       asprintf (&fmt_long, "%s:%d: error: %s: %%s\n",
 		xloc.file, xloc.line, str);
 
-      tree error_code =
-	build1 (NEGATE_EXPR, TREE_TYPE (error_var), error_var);
       print =
 	build_call_expr (builtin_decl_explicit (BUILT_IN_PRINTF), 2,
 			 build_string_literal (strlen (fmt_long) + 1,
 					       fmt_long),
-			 build_call_expr (strerror_fn, 1, error_code));
+			 build_error_message (error_var));
     }
   else
     {
@@ -480,6 +479,20 @@ build_error_statements (location_t loc, tree error_var, const char *fmt, ...)
   return stmts;
 }
 
+/* Build an error string for the StarPU return value in ERROR_VAR.  */
+
+static tree
+build_starpu_error_string (tree error_var)
+{
+  static tree strerror_fn;
+  LOOKUP_STARPU_FUNCTION (strerror_fn, "strerror");
+
+  tree error_code =
+    build1 (NEGATE_EXPR, TREE_TYPE (error_var), error_var);
+
+  return build_call_expr (strerror_fn, 1, error_code);
+}
+
 
 /* List and vector utilities, à la SRFI-1.  */
 
@@ -620,6 +633,7 @@ handle_pragma_initialize (struct cpp_reader *reader)
 		      build2 (NE_EXPR, boolean_type_node,
 			      error_var, integer_zero_node),
 		      build_error_statements (loc, error_var,
+					      build_starpu_error_string,
 					      "failed to initialize StarPU"), 
 		      NULL_TREE);
 
@@ -1140,6 +1154,113 @@ opencl_program_type (void)
   return t;
 }
 
+static tree
+opencl_kernel_type (void)
+{
+  tree t = lookup_name (get_identifier ("cl_kernel"));
+  gcc_assert (t != NULL_TREE);
+  if (TREE_CODE (t) == TYPE_DECL)
+    t = TREE_TYPE (t);
+  gcc_assert (TYPE_P (t));
+  return t;
+}
+
+static tree
+opencl_command_queue_type (void)
+{
+  tree t = lookup_name (get_identifier ("cl_command_queue"));
+  gcc_assert (t != NULL_TREE);
+  if (TREE_CODE (t) == TYPE_DECL)
+    t = TREE_TYPE (t);
+  gcc_assert (TYPE_P (t));
+  return t;
+}
+
+static tree
+opencl_event_type (void)
+{
+  tree t = lookup_name (get_identifier ("cl_event"));
+  gcc_assert (t != NULL_TREE);
+  if (TREE_CODE (t) == TYPE_DECL)
+    t = TREE_TYPE (t);
+  gcc_assert (TYPE_P (t));
+  return t;
+}
+
+/* Return an error-checking `clSetKernelArg' call for argument ARG, at
+   index IDX, of KERNEL.  */
+
+static tree
+build_opencl_set_kernel_arg_call (location_t loc, tree fn,
+				  tree kernel, unsigned int idx,
+				  tree arg)
+{
+  gcc_assert (TREE_CODE (fn) == FUNCTION_DECL
+	      && TREE_TYPE (kernel) == opencl_kernel_type ());
+
+  static tree setkernarg_fn, clstrerror_fn;
+  LOOKUP_STARPU_FUNCTION (setkernarg_fn, "clSetKernelArg");
+  LOOKUP_STARPU_FUNCTION (clstrerror_fn, "starpu_opencl_error_string");
+
+  local_define (tree, build_errorstr, (tree error_var))
+  {
+    return build_call_expr (clstrerror_fn, 1, error_var);
+  };
+
+  tree call = build_call_expr (setkernarg_fn, 4, kernel,
+			       build_int_cst (integer_type_node, idx),
+			       size_in_bytes (TREE_TYPE (arg)),
+			       build_addr (arg, fn));
+  tree error_var = build_decl (loc, VAR_DECL,
+			       create_tmp_var_name ("setkernelarg_error"),
+			       integer_type_node);
+  DECL_ARTIFICIAL (error_var) = true;
+  DECL_CONTEXT (error_var) = fn;
+
+  tree assignment = build2 (INIT_EXPR, TREE_TYPE (error_var),
+			    error_var, call);
+
+  /* Build `if (ERROR_VAR != 0) error ();'.  */
+  tree cond;
+  cond = build3 (COND_EXPR, void_type_node,
+		 build2 (NE_EXPR, boolean_type_node,
+			 error_var, integer_zero_node),
+		 build_error_statements (loc, error_var, build_errorstr,
+					 "failed to set OpenCL kernel "
+					 "argument %d", idx),
+		 NULL_TREE);
+
+  tree stmts = NULL_TREE;
+  append_to_statement_list (assignment, &stmts);
+  append_to_statement_list (cond, &stmts);
+
+  return build4 (TARGET_EXPR, void_type_node, error_var,
+		 stmts, NULL_TREE, NULL_TREE);
+}
+
+/* Return the sequence of `clSetKernelArg' calls for KERNEL.  */
+
+static tree
+build_opencl_set_kernel_arg_calls (location_t loc, tree task_impl,
+				   tree kernel)
+{
+  gcc_assert (task_implementation_p (task_impl));
+
+  size_t n;
+  tree arg, stmts = NULL_TREE;
+
+  for (arg = DECL_ARGUMENTS (task_impl), n = 0;
+       arg != NULL_TREE;
+       arg = TREE_CHAIN (arg), n++)
+    {
+      tree call = build_opencl_set_kernel_arg_call (loc, task_impl,
+						    kernel, n, arg);
+      append_to_statement_list (call, &stmts);
+    }
+
+  return stmts;
+}
+
 /* Define a body for TASK_IMPL that loads OpenCL source from FILE and calls
    KERNEL.  */
 
@@ -1151,11 +1272,21 @@ define_opencl_task_implementation (location_t loc, tree task_impl,
 	      && task_implementation_where (task_impl) == STARPU_OPENCL);
   gcc_assert (TREE_CODE (kernel) == STRING_CST);
 
+  local_define (tree, local_var, (tree type))
+  {
+    tree var = build_decl (loc, VAR_DECL,
+			   create_tmp_var_name ("opencl_var"),
+			   type);
+    DECL_ARTIFICIAL (var) = true;
+    DECL_CONTEXT (var) = task_impl;
+    return var;
+  };
+
   if (!verbose_output_p)
     /* No further warnings for this node.  */
     TREE_NO_WARNING (task_impl) = true;
 
-  static tree load_fn;
+  static tree load_fn, load_kern_fn, wid_fn, devid_fn;
 
   if (load_fn == NULL_TREE)
     {
@@ -1169,6 +1300,10 @@ define_opencl_task_implementation (location_t loc, tree task_impl,
 	}
     }
 
+  LOOKUP_STARPU_FUNCTION (load_kern_fn, "starpu_opencl_load_kernel");
+  LOOKUP_STARPU_FUNCTION (wid_fn, "starpu_worker_get_id");
+  LOOKUP_STARPU_FUNCTION (devid_fn, "starpu_worker_get_devid");
+
   if (verbose_output_p)
     inform (loc, "defining %qE, with OpenCL kernel %qs from file %qs",
 	    DECL_NAME (task_impl), TREE_STRING_POINTER (kernel), file);
@@ -1178,6 +1313,9 @@ define_opencl_task_implementation (location_t loc, tree task_impl,
 						  file, opencl_include_dirs);
   if (source_var != NULL_TREE)
     {
+      /* Give TASK_IMPL an actual argument list.  */
+      DECL_ARGUMENTS (task_impl) = build_function_arguments (task_impl);
+
       tree prog_var, prog_loaded_var;
 
       /* Global variable to hold the `starpu_opencl_program' object.  */
@@ -1219,24 +1357,59 @@ define_opencl_task_implementation (location_t loc, tree task_impl,
 
       /* Build `if (!PROG_LOADED_VAR) { ...; PROG_LOADED_VAR = true; }'.  */
 
-      tree cond = build3 (COND_EXPR, void_type_node,
-			  prog_loaded_var,
-			  NULL_TREE,
-			  load_stmts);
-
-      /* TODO: Build the kernel invocation.  */
+      tree load_cond = build3 (COND_EXPR, void_type_node,
+			       prog_loaded_var,
+			       NULL_TREE,
+			       load_stmts);
+
+      /* Local variables.  */
+      tree kernel_var, queue_var, event_var, group_size_var, ngroups_var;
+
+      kernel_var = local_var (opencl_kernel_type ());
+      queue_var = local_var (opencl_command_queue_type ());
+      event_var = local_var (opencl_event_type ());
+      group_size_var = local_var (size_type_node);
+      ngroups_var = local_var (size_type_node);
+
+      /* Build `starpu_opencl_load_kernel (...)'.
+         TODO: Check return value.  */
+      tree devid =
+	build_call_expr (devid_fn, 1, build_call_expr (wid_fn, 0));
+
+      tree load_kern = build_call_expr (load_kern_fn, 5,
+					build_addr (kernel_var, task_impl),
+					build_addr (queue_var, task_impl),
+					build_addr (prog_var, task_impl),
+					build_string_literal
+					(TREE_STRING_LENGTH (kernel) + 1,
+					 TREE_STRING_POINTER (kernel)),
+					devid);
+
+      /* TODO: `clSetKernelArg', `clEnqueueNDRangeKernel', etc.  */
+
+      /* Put it all together.  */
+      tree stmts = NULL_TREE;
+      append_to_statement_list (load_cond, &stmts);
+      append_to_statement_list (load_kern, &stmts);
+      append_to_statement_list (build_opencl_set_kernel_arg_calls (loc,
+								   task_impl,
+								   kernel_var),
+				&stmts);
+
+      /* Bind the local vars.  */
+      tree vars = chain_trees (kernel_var, queue_var, event_var,
+			       group_size_var, ngroups_var, NULL_TREE);
+      tree bind = build3 (BIND_EXPR, void_type_node, vars, stmts,
+			  build_block (vars, NULL_TREE, task_impl, NULL_TREE));
 
       TREE_USED (task_impl) = true;
       TREE_STATIC (task_impl) = true;
       DECL_EXTERNAL (task_impl) = false;
       DECL_ARTIFICIAL (task_impl) = true;
-      DECL_SAVED_TREE (task_impl) = cond;
-      DECL_INITIAL (task_impl) =
-	build_block (NULL_TREE, NULL_TREE, task_impl, NULL_TREE);
+      DECL_SAVED_TREE (task_impl) = bind;
+      DECL_INITIAL (task_impl) = BIND_EXPR_BLOCK (bind);
       DECL_RESULT (task_impl) =
 	build_decl (loc, RESULT_DECL, NULL_TREE, void_type_node);
-      DECL_ARGUMENTS (task_impl) =
-	build_function_arguments (task_impl);
 
       /* Compile TASK_IMPL.  */
       rest_of_decl_compilation (task_impl, true, 0);
@@ -2712,6 +2885,7 @@ build_pointer_lookup (tree pointer)
 		      build2 (EQ_EXPR, boolean_type_node,
 			      result_var, null_pointer_node),
 		      build_error_statements (loc, NULL_TREE,
+					      build_starpu_error_string,
 					      "attempt to use unregistered "
 					      "pointer"),
 		      NULL_TREE);
@@ -2835,6 +3009,7 @@ define_task (tree task_decl)
 		      build2 (NE_EXPR, boolean_type_node,
 			      error_var, integer_zero_node),
 		      build_error_statements (loc, error_var,
+					      build_starpu_error_string,
 					      "failed to insert task `%s'",
 					      IDENTIFIER_POINTER (name)),
 		      NULL_TREE);

gcc-plugin/tests/mocks.h

@@ -434,11 +434,19 @@ struct starpu_opencl_program
   /* Nothing.  */
 };
 
+typedef int cl_event;
+typedef int cl_kernel;
+typedef int cl_command_queue;
+
+extern cl_int clSetKernelArg (cl_kernel, cl_uint, size_t, const void *);
+
 #endif
 
 
-/* Number of `load_opencl_from_string' calls.  */
-static unsigned int load_opencl_calls;
+/* Number of `load_opencl_from_string', `load_kernel', and `clSetKernelArg'
+   calls.  */
+static unsigned int load_opencl_calls, load_opencl_kernel_calls,
+  opencl_set_kernel_arg_calls;
 
 struct load_opencl_arguments
 {
@@ -460,6 +468,72 @@ starpu_opencl_load_opencl_from_string (const char *source,
   return 0;
 }
 
+int
+starpu_opencl_load_kernel (cl_kernel *kernel,
+			   cl_command_queue *queue,
+			   struct starpu_opencl_program *programs,
+			   const char *kernel_name, int devid)
+{
+  assert (kernel != NULL && queue != NULL && programs != NULL
+	  && kernel_name != NULL && devid == -42);
+  load_opencl_kernel_calls++;
+  return 0;
+}
+
+int
+starpu_worker_get_id (void)
+{
+  return 42;
+}
+
+int
+starpu_worker_get_devid (int id)
+{
+  return -id;
+}
+
+/* Set the INDEXth argument to KERNEL to the SIZE bytes pointed to by
+   VALUE.  */
+cl_int
+clSetKernelArg (cl_kernel kernel, cl_uint index, size_t size,
+		const void *value)
+{
+  size_t n;
+  const struct insert_task_argument *arg;
+
+  for (n = 0, arg = expected_insert_task_arguments;
+       n < index;
+       n++, arg++)
+    assert (arg->pointer != NULL);
+
+  switch (arg->type)
+    {
+    case STARPU_VALUE:
+      assert (size == arg->size);
+      assert (memcmp (arg->pointer, value, size) == 0);
+      break;
+
+    case STARPU_RW:
+    case STARPU_R:
+    case STARPU_W:
+      assert (size == sizeof (void *));
+      assert (* (void **) value == arg->pointer);
+      break;
+
+    default:
+      abort ();
+    }
+
+  opencl_set_kernel_arg_calls++;
+  return 0;
+}
+
+const char *
+starpu_opencl_error_string (cl_int s)
+{
+  return "mock";
+}
+
 
 /* Initialization.  */
 

gcc-plugin/tests/opencl.c

@@ -62,6 +62,8 @@ main ()
 
   assert (tasks_submitted == 3);
   assert (load_opencl_calls == 1);
+  assert (load_opencl_kernel_calls == 3);
+  assert (opencl_set_kernel_arg_calls == 3 * 2);
 
   return EXIT_SUCCESS;
 }

include/starpu_opencl.h

@@ -31,6 +31,7 @@ extern "C"
 {
 #endif
 
+const char *starpu_opencl_error_string(cl_int status);
 void starpu_opencl_display_error(const char *func, const char *file, int line, const char *msg, cl_int status);
 #define STARPU_OPENCL_DISPLAY_ERROR(status) \
 	starpu_opencl_display_error(__starpu_func__, __FILE__, __LINE__, NULL, status)

src/core/workers.c

@@ -1136,7 +1136,7 @@ starpu_driver_init(struct starpu_driver *d)
 #endif
 #ifdef STARPU_USE_OPENCL
 	case STARPU_OPENCL_WORKER:
-		return _starpu_opencl_driver_deinit(d);
+		return _starpu_opencl_driver_init(d);
 #endif
 	case STARPU_CPU_WORKER:    /* Not supported yet */
 	case STARPU_GORDON_WORKER: /* Not supported yet */

src/drivers/opencl/driver_opencl_utils.c

@@ -330,13 +330,13 @@ int starpu_opencl_collect_stats(cl_event event STARPU_ATTRIBUTE_UNUSED)
 	return 0;
 }
 
-void starpu_opencl_display_error(const char *func, const char *file, int line, const char* msg, cl_int status)
+const char *starpu_opencl_error_string(cl_int status)
 {
 	const char *errormsg;
 	switch (status)
 	{
 	case CL_SUCCESS:
-		errormsg = "success";
+		errormsg = "Success";
 		break;
 	case CL_DEVICE_NOT_FOUND:
 		errormsg = "Device not found";
@@ -479,13 +479,16 @@ void starpu_opencl_display_error(const char *func, const char *file, int line, c
 		break;
 #endif
 	default:
-		errormsg = "unknown error";
+		errormsg = "unknown OpenCL error";
 		break;
 	}
-	if (msg)
-		printf("oops in %s (%s:%d) (%s) ... <%s> (%d) \n", func, file, line, msg, errormsg, status);
-	else
-		printf("oops in %s (%s:%d) ... <%s> (%d) \n", func, file, line, errormsg, status);
+	return errormsg;
+}
+
+void starpu_opencl_display_error(const char *func, const char *file, int line, const char* msg, cl_int status)
+{
+	printf("oops in %s (%s:%d) (%s) ... <%s> (%d) \n", func, file, line, msg,
+	       starpu_opencl_error_string (status), status);
 
 }