merge trunk@6520:6550

.gitignore

@@ -286,3 +286,4 @@ starpu.log
 /tools/starpu_workers_activity
 /tests/datawizard/interfaces/copy_interfaces
 /gcc-plugin/tests/release
+/gcc-plugin/tests/opencl

build-aux/compile

@@ -1,9 +1,9 @@
 #! /bin/sh
-# Wrapper for compilers which do not understand `-c -o'.
+# Wrapper for compilers which do not understand '-c -o'.
 
-scriptversion=2005-05-14.22
+scriptversion=2012-03-05.13; # UTC
 
-# Copyright (C) 1999, 2000, 2003, 2004, 2005 Free Software Foundation, Inc.
+# Copyright (C) 1999-2012 Free Software Foundation, Inc.
 # Written by Tom Tromey <tromey@cygnus.com>.
 #
 # This program is free software; you can redistribute it and/or modify
@@ -17,8 +17,7 @@ scriptversion=2005-05-14.22
 # GNU General Public License for more details.
 #
 # You should have received a copy of the GNU General Public License
-# along with this program; if not, write to the Free Software
-# Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
+# along with this program.  If not, see <http://www.gnu.org/licenses/>.
 
 # As a special exception to the GNU General Public License, if you
 # distribute this file as part of a program that contains a
@@ -29,21 +28,219 @@ scriptversion=2005-05-14.22
 # bugs to <bug-automake@gnu.org> or send patches to
 # <automake-patches@gnu.org>.
 
+nl='
+'
+
+# We need space, tab and new line, in precisely that order.  Quoting is
+# there to prevent tools from complaining about whitespace usage.
+IFS=" ""	$nl"
+
+file_conv=
+
+# func_file_conv build_file lazy
+# Convert a $build file to $host form and store it in $file
+# Currently only supports Windows hosts. If the determined conversion
+# type is listed in (the comma separated) LAZY, no conversion will
+# take place.
+func_file_conv ()
+{
+  file=$1
+  case $file in
+    / | /[!/]*) # absolute file, and not a UNC file
+      if test -z "$file_conv"; then
+	# lazily determine how to convert abs files
+	case `uname -s` in
+	  MINGW*)
+	    file_conv=mingw
+	    ;;
+	  CYGWIN*)
+	    file_conv=cygwin
+	    ;;
+	  *)
+	    file_conv=wine
+	    ;;
+	esac
+      fi
+      case $file_conv/,$2, in
+	*,$file_conv,*)
+	  ;;
+	mingw/*)
+	  file=`cmd //C echo "$file " | sed -e 's/"\(.*\) " *$/\1/'`
+	  ;;
+	cygwin/*)
+	  file=`cygpath -m "$file" || echo "$file"`
+	  ;;
+	wine/*)
+	  file=`winepath -w "$file" || echo "$file"`
+	  ;;
+      esac
+      ;;
+  esac
+}
+
+# func_cl_dashL linkdir
+# Make cl look for libraries in LINKDIR
+func_cl_dashL ()
+{
+  func_file_conv "$1"
+  if test -z "$lib_path"; then
+    lib_path=$file
+  else
+    lib_path="$lib_path;$file"
+  fi
+  linker_opts="$linker_opts -LIBPATH:$file"
+}
+
+# func_cl_dashl library
+# Do a library search-path lookup for cl
+func_cl_dashl ()
+{
+  lib=$1
+  found=no
+  save_IFS=$IFS
+  IFS=';'
+  for dir in $lib_path $LIB
+  do
+    IFS=$save_IFS
+    if $shared && test -f "$dir/$lib.dll.lib"; then
+      found=yes
+      lib=$dir/$lib.dll.lib
+      break
+    fi
+    if test -f "$dir/$lib.lib"; then
+      found=yes
+      lib=$dir/$lib.lib
+      break
+    fi
+  done
+  IFS=$save_IFS
+
+  if test "$found" != yes; then
+    lib=$lib.lib
+  fi
+}
+
+# func_cl_wrapper cl arg...
+# Adjust compile command to suit cl
+func_cl_wrapper ()
+{
+  # Assume a capable shell
+  lib_path=
+  shared=:
+  linker_opts=
+  for arg
+  do
+    if test -n "$eat"; then
+      eat=
+    else
+      case $1 in
+	-o)
+	  # configure might choose to run compile as 'compile cc -o foo foo.c'.
+	  eat=1
+	  case $2 in
+	    *.o | *.[oO][bB][jJ])
+	      func_file_conv "$2"
+	      set x "$@" -Fo"$file"
+	      shift
+	      ;;
+	    *)
+	      func_file_conv "$2"
+	      set x "$@" -Fe"$file"
+	      shift
+	      ;;
+	  esac
+	  ;;
+	-I)
+	  eat=1
+	  func_file_conv "$2" mingw
+	  set x "$@" -I"$file"
+	  shift
+	  ;;
+	-I*)
+	  func_file_conv "${1#-I}" mingw
+	  set x "$@" -I"$file"
+	  shift
+	  ;;
+	-l)
+	  eat=1
+	  func_cl_dashl "$2"
+	  set x "$@" "$lib"
+	  shift
+	  ;;
+	-l*)
+	  func_cl_dashl "${1#-l}"
+	  set x "$@" "$lib"
+	  shift
+	  ;;
+	-L)
+	  eat=1
+	  func_cl_dashL "$2"
+	  ;;
+	-L*)
+	  func_cl_dashL "${1#-L}"
+	  ;;
+	-static)
+	  shared=false
+	  ;;
+	-Wl,*)
+	  arg=${1#-Wl,}
+	  save_ifs="$IFS"; IFS=','
+	  for flag in $arg; do
+	    IFS="$save_ifs"
+	    linker_opts="$linker_opts $flag"
+	  done
+	  IFS="$save_ifs"
+	  ;;
+	-Xlinker)
+	  eat=1
+	  linker_opts="$linker_opts $2"
+	  ;;
+	-*)
+	  set x "$@" "$1"
+	  shift
+	  ;;
+	*.cc | *.CC | *.cxx | *.CXX | *.[cC]++)
+	  func_file_conv "$1"
+	  set x "$@" -Tp"$file"
+	  shift
+	  ;;
+	*.c | *.cpp | *.CPP | *.lib | *.LIB | *.Lib | *.OBJ | *.obj | *.[oO])
+	  func_file_conv "$1" mingw
+	  set x "$@" "$file"
+	  shift
+	  ;;
+	*)
+	  set x "$@" "$1"
+	  shift
+	  ;;
+      esac
+    fi
+    shift
+  done
+  if test -n "$linker_opts"; then
+    linker_opts="-link$linker_opts"
+  fi
+  exec "$@" $linker_opts
+  exit 1
+}
+
+eat=
+
 case $1 in
   '')
-     echo "$0: No command.  Try \`$0 --help' for more information." 1>&2
+     echo "$0: No command.  Try '$0 --help' for more information." 1>&2
      exit 1;
      ;;
   -h | --h*)
     cat <<\EOF
 Usage: compile [--help] [--version] PROGRAM [ARGS]
 
-Wrapper for compilers which do not understand `-c -o'.
-Remove `-o dest.o' from ARGS, run PROGRAM with the remaining
+Wrapper for compilers which do not understand '-c -o'.
+Remove '-o dest.o' from ARGS, run PROGRAM with the remaining
 arguments, and rename the output as expected.
 
 If you are trying to build a whole package this is not the
-right script to run: please start by reading the file `INSTALL'.
+right script to run: please start by reading the file 'INSTALL'.
 
 Report bugs to <bug-automake@gnu.org>.
 EOF
@@ -53,11 +250,13 @@ EOF
     echo "compile $scriptversion"
     exit $?
     ;;
+  cl | *[/\\]cl | cl.exe | *[/\\]cl.exe )
+    func_cl_wrapper "$@"      # Doesn't return...
+    ;;
 esac
 
 ofile=
 cfile=
-eat=
 
 for arg
 do
@@ -66,8 +265,8 @@ do
   else
     case $1 in
       -o)
-	# configure might choose to run compile as `compile cc -o foo foo.c'.
-	# So we strip `-o arg' only if arg is an object.
+	# configure might choose to run compile as 'compile cc -o foo foo.c'.
+	# So we strip '-o arg' only if arg is an object.
 	eat=1
 	case $2 in
 	  *.o | *.obj)
@@ -94,22 +293,22 @@ do
 done
 
 if test -z "$ofile" || test -z "$cfile"; then
-  # If no `-o' option was seen then we might have been invoked from a
+  # If no '-o' option was seen then we might have been invoked from a
   # pattern rule where we don't need one.  That is ok -- this is a
   # normal compilation that the losing compiler can handle.  If no
-  # `.c' file was seen then we are probably linking.  That is also
+  # '.c' file was seen then we are probably linking.  That is also
   # ok.
   exec "$@"
 fi
 
 # Name of file we expect compiler to create.
-cofile=`echo "$cfile" | sed -e 's|^.*/||' -e 's/\.c$/.o/'`
+cofile=`echo "$cfile" | sed 's|^.*[\\/]||; s|^[a-zA-Z]:||; s/\.c$/.o/'`
 
 # Create the lock directory.
-# Note: use `[/.-]' here to ensure that we don't use the same name
+# Note: use '[/\\:.-]' here to ensure that we don't use the same name
 # that we are using for the .o file.  Also, base the name on the expected
 # object file name, since that is what matters with a parallel build.
-lockdir=`echo "$cofile" | sed -e 's|[/.-]|_|g'`.d
+lockdir=`echo "$cofile" | sed -e 's|[/\\:.-]|_|g'`.d
 while true; do
   if mkdir "$lockdir" >/dev/null 2>&1; then
     break
@@ -124,9 +323,9 @@ trap "rmdir '$lockdir'; exit 1" 1 2 15
 ret=$?
 
 if test -f "$cofile"; then
-  mv "$cofile" "$ofile"
+  test "$cofile" = "$ofile" || mv "$cofile" "$ofile"
 elif test -f "${cofile}bj"; then
-  mv "${cofile}bj" "$ofile"
+  test "${cofile}bj" = "$ofile" || mv "${cofile}bj" "$ofile"
 fi
 
 rmdir "$lockdir"
@@ -138,5 +337,6 @@ exit $ret
 # eval: (add-hook 'write-file-hooks 'time-stamp)
 # time-stamp-start: "scriptversion="
 # time-stamp-format: "%:y-%02m-%02d.%02H"
-# time-stamp-end: "$"
+# time-stamp-time-zone: "UTC"
+# time-stamp-end: "; # UTC"
 # End:

configure.ac

@@ -59,7 +59,9 @@ m4_ifdef([AM_SILENT_RULES],
 
 AC_PREREQ(2.60)
 
+m4_ifdef([AM_PROG_AR], [AM_PROG_AR])
 AC_PROG_CC
+AM_PROG_CC_C_O
 AC_PROG_CXX
 AC_PROG_CPP
 AC_PROG_SED
@@ -132,9 +134,9 @@ else
 fi
 
 AC_COMPILE_IFELSE(
-  AC_LANG_PROGRAM([[
+  [AC_LANG_PROGRAM([[
     #include <pthread.h>
-  ]], [[ pthread_t t; pthread_create(&t, NULL, NULL, NULL); ]]),,
+  ]], [[ pthread_t t; pthread_create(&t, NULL, NULL, NULL); ]])],,
   AC_MSG_ERROR([pthread_create unavailable]))
 AC_SEARCH_LIBS([sqrt],[m],,AC_MSG_ERROR([math library unavailable]))
 AC_HAVE_LIBRARY([ws2_32])
@@ -1578,11 +1580,11 @@ if test "$enable_cuda" = "yes" -a "$ICC" != ""; then
    OLD_CFLAGS="$CFLAGS"
    CFLAGS="-I$PWD/include -I$srcdir/include"
    AC_COMPILE_IFELSE(
-       AC_LANG_PROGRAM(
+       [AC_LANG_PROGRAM(
 	   [[#include <cuda.h>
 	   #include <starpu.h>]],
 	   [[]]
-	   ),
+	   )],
        AC_MSG_RESULT(yes),
        [ICC=""
            AC_MSG_RESULT(no)]

doc/Makefile.am

@@ -60,7 +60,7 @@ uninstall-local:
 #	vector_scal_c.texi vector_scal_cuda.texi vector_scal_opencl.texi vector_scal_opencl_codelet.texi
 
 # Rule to update documentation on web server. Should only be used locally.
-PUBLISHHOST	:= sync
+PUBLISHHOST	= sync
 update-web: starpu.html
 	sed -i 's/gcc\.html#Attribute-Syntax/http:\/\/gcc.gnu.org\/onlinedocs\/gcc\/Attribute-Syntax.html#Attribute-Syntax/' starpu.html
 	scp starpu.pdf starpu.html $(PUBLISHHOST):/web/runtime/html/StarPU

doc/chapters/advanced-examples.texi

@@ -932,7 +932,7 @@ be enabled by using the @code{cuda_opengl_interoperability} field of the
 @code{starpu_conf} structure, and the driver loop has to be run by
 the application, by using the @code{not_launched_drivers} field of
 @code{starpu_conf} to prevent StarPU from running it in a separate thread, and
-by using @code{starpu_run_driver} to run the loop. The @code{gl_interop} example
+by using @code{starpu_driver_run} to run the loop. The @code{gl_interop} example
 shows how it articulates in a simple case, where rendering is done in task
 callbacks. TODO: provide glutIdleFunc alternative.
 

examples/gl_interop/gl_interop.c

@@ -120,7 +120,7 @@ int main(int argc, char **argv)
 
 	/* And run the driver, which will run the task */
 	printf("running the driver\n");
-	starpu_run_driver(&drivers[0]);
+	starpu_driver_run(&drivers[0]);
 	printf("finished running the driver\n");
 
 	starpu_shutdown();

gcc-plugin/src/c-expr.y

@@ -89,6 +89,24 @@
     sorry ("struct field access not implemented yet"); /* XXX */
     return error_mark_node;
   }
+
+  /* Interpret the string beneath CST, and return a new string constant.  */
+  static tree
+  interpret_string (const_tree cst)
+  {
+    gcc_assert (TREE_CODE (cst) == STRING_CST);
+
+    cpp_string input, interpreted;
+    input.text = (unsigned char *) TREE_STRING_POINTER (cst);
+    input.len = TREE_STRING_LENGTH (cst);
+
+    bool success;
+    success = cpp_interpret_string (parse_in, &input, 1, &interpreted,
+				    CPP_STRING);
+    gcc_assert (success);
+
+    return build_string (interpreted.len, (char *) interpreted.text);
+  }
 %}
 
 %code {
@@ -132,6 +150,8 @@
   yylex (YYSTYPE *lvalp)
   {
     int ret;
+    enum cpp_ttype type;
+    location_t loc;
 
 #ifdef __cplusplus
     if (cpplib_bison_token_map[CPP_NAME] != YCPP_NAME)
@@ -143,11 +163,30 @@
       }
 #endif
 
-    ret = pragma_lex (lvalp);
-    if (ret < sizeof cpplib_bison_token_map / sizeof cpplib_bison_token_map[0])
-      ret = cpplib_bison_token_map[ret];
-    else
+    /* First check whether EOL is reached, because the EOL token needs to be
+       left to the C parser.  */
+    type = cpp_peek_token (parse_in, 0)->type;
+    if (type == CPP_PRAGMA_EOL)
       ret = -1;
+    else
+      {
+	/* Tell the lexer to not concatenate adjacent strings like cpp and
+	   `pragma_lex' normally do, because we want to be able to
+	   distinguish adjacent STRING_CST.  */
+	type = c_lex_with_flags (lvalp, &loc, NULL, C_LEX_STRING_NO_JOIN);
+
+	if (type == CPP_STRING)
+	  /* XXX: When using `C_LEX_STRING_NO_JOIN', `c_lex_with_flags'
+	     doesn't call `cpp_interpret_string', leaving us with an
+	     uninterpreted string (with quotes, etc.)  This hack works around
+	     that.  */
+	  *lvalp = interpret_string (*lvalp);
+
+	if (type < sizeof cpplib_bison_token_map / sizeof cpplib_bison_token_map[0])
+	  ret = cpplib_bison_token_map[type];
+	else
+	  ret = -1;
+      }
 
     return ret;
   }

gcc-plugin/src/starpu.c

@@ -54,6 +54,7 @@
 #include <toplev.h>
 
 #include <stdio.h>
+#include <sys/mman.h>
 
 /* Don't include the dreaded proprietary headers that we don't need anyway.
    In particular, this waives the obligation to reproduce their silly
@@ -115,6 +116,9 @@ static const char plugin_name[] = "starpu";
 /* Whether to enable verbose output.  */
 static bool verbose_output_p = false;
 
+/* Search path for OpenCL source files, for the `opencl' pragma.  */
+static tree opencl_include_dirs = NULL_TREE;
+
 /* Names of public attributes.  */
 static const char task_attribute_name[] = "task";
 static const char task_implementation_attribute_name[] = "task_implementation";
@@ -132,21 +136,41 @@ static const char heap_allocated_orig_type_attribute_name[] =
 
 /* Names of data structures defined in <starpu.h>.  */
 static const char codelet_struct_tag[] = "starpu_codelet";
+static const char opencl_program_struct_tag[] = "starpu_opencl_program";
 
 /* Cached function declarations.  */
 static tree unpack_fn, data_lookup_fn;
 
+/* Targets supported by GCC-StarPU.  */
+static int supported_targets = 0
+#ifdef STARPU_USE_CPU
+    | STARPU_CPU
+#endif
+#ifdef STARPU_USE_CUDA
+    | STARPU_CUDA
+#endif
+#ifdef STARPU_USE_OPENCL
+    | STARPU_OPENCL
+#endif
+#ifdef STARPU_USE_GORDON
+    | STARPU_GORDON
+#endif
+    ;
+
 
 /* Forward declarations.  */
 
+static tree build_function_arguments (tree fn);
 static tree build_codelet_declaration (tree task_decl);
 static tree build_cpu_codelet_identifier (const_tree task);
 static void define_task (tree task_decl);
 static tree build_pointer_lookup (tree pointer);
+static tree type_decl_for_struct_tag (const char *tag);
 
 static bool task_p (const_tree decl);
 static bool task_implementation_p (const_tree decl);
 static tree task_implementation_task (const_tree task_impl);
+static int task_implementation_where (const_tree task_impl);
 static bool implicit_cpu_task_implementation_p (const_tree fn);
 
 
@@ -154,6 +178,8 @@ static int task_implementation_target_to_int (const_tree target);
 
 static bool heap_allocated_p (const_tree var_decl);
 
+static tree declare_codelet (tree task_decl);
+
 
 /* Lookup the StarPU function NAME in the global scope and store the result
    in VAR (this can't be done from `lower_starpu'.)  */
@@ -769,6 +795,13 @@ handle_pragma_register (struct cpp_reader *reader)
   else
     ptr_type = TREE_TYPE (ptr);
 
+  if (ptr_type == NULL_TREE)
+    {
+      /* PTR is a type-less thing, such as a STRING_CST.  */
+      error_at (loc, "invalid %<register%> argument");
+      return;
+    }
+
   if (!POINTER_TYPE_P (ptr_type)
       && TREE_CODE (ptr_type) != ARRAY_TYPE)
     {
@@ -988,6 +1021,315 @@ handle_pragma_unregister (struct cpp_reader *reader)
   add_stmt (build_data_unregister_call (loc, var));
 }
 
+/* Return a private global string literal VAR_DECL, whose contents are the
+   LEN bytes at CONTENTS.  */
+
+static tree
+build_string_variable (location_t loc, const char *name_seed,
+		       const char *contents, size_t len)
+{
+  tree decl;
+
+  decl = build_decl (loc, VAR_DECL, create_tmp_var_name (name_seed),
+		     string_type_node);
+  TREE_PUBLIC (decl) = false;
+  TREE_STATIC (decl) = true;
+  TREE_USED (decl) = true;
+
+  DECL_INITIAL (decl) =				  /* XXX: off-by-one? */
+    build_string_literal (len + 1, contents);
+
+  DECL_ARTIFICIAL (decl) = true;
+
+  return decl;
+}
+
+/* Return a VAR_DECL for a string variable containing the contents of FILE,
+   which is looked for in each of the directories listed in SEARCH_PATH.  If
+   FILE could not be found, return NULL_TREE.  */
+
+static tree
+build_variable_from_file_contents (location_t loc,
+				   const char *name_seed,
+				   const char *file,
+				   const_tree search_path)
+{
+  gcc_assert (search_path != NULL_TREE
+	      && TREE_CODE (search_path) == STRING_CST);
+
+  int err, dir_fd;
+  struct stat st;
+  const_tree dirs;
+  tree var = NULL_TREE;
+
+  /* Look for FILE in each directory in SEARCH_PATH, and pick the first one
+     that matches.  */
+  for (err = ENOENT, dir_fd = -1, dirs = opencl_include_dirs;
+       (err != 0 || err == ENOENT) && dirs != NULL_TREE;
+       dirs = TREE_CHAIN (dirs))
+    {
+      dir_fd = open (TREE_STRING_POINTER (dirs),
+		     O_DIRECTORY | O_RDONLY);
+      if (dir_fd < 0)
+	err = ENOENT;
+      else
+	{
+	  err = fstatat (dir_fd, file, &st, 0);
+	  if (err != 0)
+	    close (dir_fd);
+	}
+    }
+
+  if (err != 0 || dir_fd < 0)
+    error_at (loc, "failed to access %qs: %m", file);
+  else if (st.st_size == 0)
+    {
+      error_at (loc, "source file %qs is empty", file);
+      close (dir_fd);
+    }
+  else
+    {
+      if (verbose_output_p)
+	inform (loc, "found file %qs in %qs",
+		file, TREE_STRING_POINTER (dirs));
+
+      int fd;
+
+      fd = openat (dir_fd, file, O_RDONLY);
+      close (dir_fd);
+
+      if (fd < 0)
+	error_at (loc, "failed to open %qs: %m", file);
+      else
+	{
+	  void *contents;
+
+	  contents = mmap (NULL, st.st_size, PROT_READ, MAP_SHARED, fd, 0);
+	  if (contents == NULL)
+	    error_at (loc, "failed to map contents of %qs: %m", file);
+	  else
+	    {
+	      var = build_string_variable (loc, name_seed,
+					   (char *) contents, st.st_size);
+	      pushdecl (var);
+	      munmap (contents, st.st_size);
+	    }
+
+	  close (fd);
+	}
+    }
+
+  return var;
+}
+
+/* Return the type corresponding to OPENCL_PROGRAM_STRUCT_TAG.  */
+
+static tree
+opencl_program_type (void)
+{
+  tree t = TREE_TYPE (type_decl_for_struct_tag (opencl_program_struct_tag));
+
+  if (TYPE_SIZE (t) == NULL_TREE)
+    {
+      /* Incomplete type definition, for instance because <starpu_opencl.h>
+	 wasn't included.  */
+      error_at (UNKNOWN_LOCATION, "StarPU OpenCL support is lacking");
+      t = error_mark_node;
+    }
+
+  return t;
+}
+
+/* Define a body for TASK_IMPL that loads OpenCL source from FILE and calls
+   KERNEL.  */
+
+static void
+define_opencl_task_implementation (location_t loc, tree task_impl,
+				   const char *file, const_tree kernel)
+{
+  gcc_assert (task_implementation_p (task_impl)
+	      && task_implementation_where (task_impl) == STARPU_OPENCL);
+  gcc_assert (TREE_CODE (kernel) == STRING_CST);
+
+  if (!verbose_output_p)
+    /* No further warnings for this node.  */
+    TREE_NO_WARNING (task_impl) = true;
+
+  static tree load_fn;
+
+  if (load_fn == NULL_TREE)
+    {
+      load_fn =
+	lookup_name (get_identifier ("starpu_opencl_load_opencl_from_string"));
+      if (load_fn == NULL_TREE)
+	{
+	  inform (loc, "no OpenCL support, task implementation %qE "
+		  "not generated", DECL_NAME (task_impl));
+	  return;
+	}
+    }
+
+  if (verbose_output_p)
+    inform (loc, "defining %qE, with OpenCL kernel %qs from file %qs",
+	    DECL_NAME (task_impl), TREE_STRING_POINTER (kernel), file);
+
+  tree source_var;
+  source_var = build_variable_from_file_contents (loc, "opencl_source",
+						  file, opencl_include_dirs);
+  if (source_var != NULL_TREE)
+    {
+      tree prog_var, prog_loaded_var;
+
+      /* Global variable to hold the `starpu_opencl_program' object.  */
+
+      prog_var = build_decl (loc, VAR_DECL,
+			     create_tmp_var_name ("opencl_program"),
+			     opencl_program_type ());
+      TREE_PUBLIC (prog_var) = false;
+      TREE_STATIC (prog_var) = true;
+      TREE_USED (prog_var) = true;
+      DECL_ARTIFICIAL (prog_var) = true;
+      pushdecl (prog_var);
+
+      /* Global variable indicating whether the program has already been
+	 loaded.  */
+
+      prog_loaded_var = build_decl (loc, VAR_DECL,
+				    create_tmp_var_name ("opencl_prog_loaded"),
+				    boolean_type_node);
+      TREE_PUBLIC (prog_loaded_var) = false;
+      TREE_STATIC (prog_loaded_var) = true;
+      TREE_USED (prog_loaded_var) = true;
+      DECL_ARTIFICIAL (prog_loaded_var) = true;
+      DECL_INITIAL (prog_loaded_var) = build_zero_cst (boolean_type_node);
+      pushdecl (prog_loaded_var);
+
+      /* Build `starpu_opencl_load_opencl_from_string (SOURCE_VAR,
+	                                               &PROG_VAR, "")'.  */
+      tree load = build_call_expr (load_fn, 3, source_var,
+				   build_addr (prog_var, task_impl),
+				   build_string_literal (1, ""));
+
+      tree load_stmts = NULL_TREE;
+      append_to_statement_list (load, &load_stmts);
+      append_to_statement_list (build2 (MODIFY_EXPR, boolean_type_node,
+					prog_loaded_var,
+					build_int_cst (boolean_type_node, 1)),
+				&load_stmts);
+
+      /* 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_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_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);
+      allocate_struct_function (task_impl, false);
+      cgraph_finalize_function (task_impl, false);
+      cgraph_mark_needed_node (cgraph_get_node (task_impl));
+
+      /* Generate a wrapper for TASK_IMPL, and possibly the body of its task.
+	 This needs to be done explicitly here, because otherwise
+	 `handle_pre_genericize' would never see TASK_IMPL's task.  */
+      tree task = task_implementation_task (task_impl);
+      if (!TREE_STATIC (task))
+	{
+	  declare_codelet (task);
+	  define_task (task);
+
+	  /* Compile TASK's body.  */
+	  rest_of_decl_compilation (task, true, 0);
+	  allocate_struct_function (task, false);
+	  cgraph_finalize_function (task, false);
+	  cgraph_mark_needed_node (cgraph_get_node (task));
+	}
+    }
+  else
+    DECL_SAVED_TREE (task_impl) = error_mark_node;
+
+  return;
+}
+
+/* Handle the `opencl' pragma, which defines an OpenCL task
+   implementation.  */
+
+static void
+handle_pragma_opencl (struct cpp_reader *reader)
+{
+  tree args;
+  location_t loc;
+
+  loc = cpp_peek_token (reader, 0)->src_loc;
+
+  if (current_function_decl != NULL_TREE)
+    {
+      error_at (loc, "%<starpu opencl%> pragma can only be used "
+		"at the top-level");
+      return;
+    }
+
+  args = read_pragma_expressions ("opencl", loc);
+  if (args == NULL_TREE)
+    return;
+
+  /* TODO: Add "group size" and "number of groups" arguments.  */
+  if (list_length (args) < 3)
+    {
+      error_at (loc, "wrong number of arguments for %<starpu opencl%> pragma");
+      return;
+    }
+
+  if (task_implementation_p (TREE_VALUE (args)))
+    {
+      tree task_impl = TREE_VALUE (args);
+      if (task_implementation_where (task_impl) == STARPU_OPENCL)
+  	{
+  	  args = TREE_CHAIN (args);
+  	  if (TREE_CODE (TREE_VALUE (args)) == STRING_CST)
+  	    {
+  	      tree file = TREE_VALUE (args);
+  	      args = TREE_CHAIN (args);
+  	      if (TREE_CODE (TREE_VALUE (args)) == STRING_CST)
+  		{
+  		  tree kernel = TREE_VALUE (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");
+  		}
+  	      else
+  		error_at (loc, "%<kernel%> argument must be a string constant");
+	    }
+	  else
+	    error_at (loc, "%<file%> argument must be a string constant");
+	}
+      else
+	error_at (loc, "%qE is not an OpenCL task implementation",
+		  DECL_NAME (task_impl));
+    }
+  else
+    error_at (loc, "%qE is not a task implementation", TREE_VALUE (args));
+}
+
 /* Handle the `debug_tree' pragma (for debugging purposes.)  */
 
 static void
@@ -1017,6 +1359,44 @@ handle_pragma_debug_tree (struct cpp_reader *reader)
   printf ("\n");
 }
 
+/* Handle the `#pragma starpu add_target TARGET', which tells GCC-StarPU to
+   consider TARGET ("cpu", "opencl", etc.) as supported.  This pragma is
+   undocumented and only meant to be used for testing purposes.  */
+
+static void
+handle_pragma_add_target (struct cpp_reader *reader)
+{
+  tree args, obj;
+  location_t loc;
+
+  loc = cpp_peek_token (reader, 0)->src_loc;
+
+  args = read_pragma_expressions ("add_target", loc);
+  if (args == NULL_TREE)
+    /* Parse error, presumably already handled by the parser.  */
+    return;
+
+  obj = TREE_VALUE (args);
+  args = TREE_CHAIN (args);
+
+  if (obj == error_mark_node)
+    return;
+
+  if (args != NULL_TREE)
+    warning_at (loc, 0, "extraneous arguments ignored");
+
+  if (TREE_CODE (obj) == STRING_CST)
+    {
+      int new_target = task_implementation_target_to_int (obj);
+      if (obj == 0)
+	error_at (loc, "unsupported target %qE", obj);
+      else
+	supported_targets |= new_target;
+    }
+  else
+    error_at (loc, "expecting string literal");
+}
+
 static void
 register_pragmas (void *gcc_data, void *user_data)
 {
@@ -1024,6 +1404,8 @@ register_pragmas (void *gcc_data, void *user_data)
 		     handle_pragma_hello);
   c_register_pragma (STARPU_PRAGMA_NAME_SPACE, "debug_tree",
 		     handle_pragma_debug_tree);
+  c_register_pragma (STARPU_PRAGMA_NAME_SPACE, "add_target",
+		     handle_pragma_add_target);
 
   c_register_pragma_with_expansion (STARPU_PRAGMA_NAME_SPACE, "initialize",
 				    handle_pragma_initialize);
@@ -1037,6 +1419,8 @@ register_pragmas (void *gcc_data, void *user_data)
 				    handle_pragma_release);
   c_register_pragma_with_expansion (STARPU_PRAGMA_NAME_SPACE, "unregister",
 				    handle_pragma_unregister);
+  c_register_pragma_with_expansion (STARPU_PRAGMA_NAME_SPACE, "opencl",
+				    handle_pragma_opencl);
   c_register_pragma (STARPU_PRAGMA_NAME_SPACE, "shutdown",
 		     handle_pragma_shutdown);
 }
@@ -1935,6 +2319,23 @@ build_codelet_identifier (tree task_decl)
   return get_identifier (cl_name);
 }
 
+/* Return a TYPE_DECL for the RECORD_TYPE with tag name TAG.  */
+
+static tree
+type_decl_for_struct_tag (const char *tag)
+{
+  tree type_decl = xref_tag (RECORD_TYPE, get_identifier (tag));
+  gcc_assert (type_decl != NULL_TREE
+	      && TREE_CODE (type_decl) == RECORD_TYPE);
+
+  /* `build_decl' expects a TYPE_DECL, so give it what it wants.  */
+
+  type_decl = TYPE_STUB_DECL (type_decl);
+  gcc_assert (type_decl != NULL && TREE_CODE (type_decl) == TYPE_DECL);
+
+  return type_decl;
+}
+
 static tree
 codelet_type (void)
 {
@@ -1943,19 +2344,9 @@ codelet_type (void)
   static tree type_decl = NULL_TREE;
 
   if (type_decl == NULL_TREE)
-    {
-      /* Lookup the `struct starpu_codelet' struct type.  This should succeed since
-	 we push <starpu.h> early on.  */
-
-      type_decl = xref_tag (RECORD_TYPE, get_identifier (codelet_struct_tag));
-      gcc_assert (type_decl != NULL_TREE
-		  && TREE_CODE (type_decl) == RECORD_TYPE);
-
-      /* `build_decl' expects a TYPE_DECL, so give it what it wants.  */
-
-      type_decl = TYPE_STUB_DECL (type_decl);
-      gcc_assert (type_decl != NULL && TREE_CODE (type_decl) == TYPE_DECL);
-    }
+    /* Lookup the `struct starpu_codelet' struct type.  This should succeed since
+       we push <starpu.h> early on.  */
+    type_decl = type_decl_for_struct_tag (codelet_struct_tag);
 
   return TREE_TYPE (type_decl);
 }
@@ -2266,30 +2657,11 @@ handle_pre_genericize (void *gcc_data, void *user_data)
 	  /* TASK lacks a body.  Declare its codelet, intantiate its codelet
 	     wrappers, and its body in this compilation unit.  */
 
-	  local_define (tree, build_parameter, (const_tree lst))
-	  {
-	    tree param, type;
-
-	    type = TREE_VALUE (lst);
-	    param = build_decl (DECL_SOURCE_LOCATION (task), PARM_DECL,
-				create_tmp_var_name ("parameter"),
-				type);
-	    DECL_ARG_TYPE (param) = type;
-	    DECL_CONTEXT (param) = task;
-
-	    return param;
-	  };
 
 	  /* Declare TASK's codelet.  It cannot be defined yet because the
 	     complete list of tasks isn't available at this point.  */
 	  declare_codelet (task);
 
-	  /* Set the task's parameter list.  */
-	  DECL_ARGUMENTS (task) =
-	    map (build_parameter,
-		 list_remove (void_type_p,
-			      TYPE_ARG_TYPES (TREE_TYPE (task))));
-
 	  /* Build its body.  */
 	  current_function_decl = task;
 	  define_task (task);
@@ -2352,11 +2724,42 @@ build_pointer_lookup (tree pointer)
   return build4 (TARGET_EXPR, ptr_type_node, result_var, stmts, NULL_TREE, NULL_TREE);
 }
 
+/* Return a fresh argument list for FN.  */
+
+static tree
+build_function_arguments (tree fn)
+{
+  gcc_assert (TREE_CODE (fn) == FUNCTION_DECL
+	      && DECL_ARGUMENTS (fn) == NULL_TREE);
+
+  local_define (tree, build_argument, (const_tree lst))
+    {
+      tree param, type;
+
+      type = TREE_VALUE (lst);
+      param = build_decl (DECL_SOURCE_LOCATION (fn), PARM_DECL,
+			  create_tmp_var_name ("argument"),
+			  type);
+      DECL_ARG_TYPE (param) = type;
+      DECL_CONTEXT (param) = fn;
+
+      return param;
+    };
+
+  return map (build_argument,
+	      list_remove (void_type_p,
+			   TYPE_ARG_TYPES (TREE_TYPE (fn))));
+}
+
+
 /* Build the body of TASK_DECL, which will call `starpu_insert_task'.  */
 
 static void
 define_task (tree task_decl)
 {
+  /* First of all, give TASK_DECL an argument list.  */
+  DECL_ARGUMENTS (task_decl) = build_function_arguments (task_decl);
+
   VEC(tree, gc) *args = NULL;
   location_t loc = DECL_SOURCE_LOCATION (task_decl);
   tree p, params = DECL_ARGUMENTS (task_decl);
@@ -2463,27 +2866,12 @@ validate_task (tree task)
 {
   gcc_assert (task_p (task));
 
-  static const int supported = 0
-#ifdef STARPU_USE_CPU
-    | STARPU_CPU
-#endif
-#ifdef STARPU_USE_CUDA
-    | STARPU_CUDA
-#endif
-#ifdef STARPU_USE_OPENCL
-    | STARPU_OPENCL
-#endif
-#ifdef STARPU_USE_GORDON
-    | STARPU_GORDON
-#endif
-    ;
-
   int where = task_where (task);
 
   /* If TASK has no implementations, things will barf elsewhere anyway.  */
 
   if (task_implementation_list (task) != NULL_TREE)
-    if ((where & supported) == 0)
+    if ((where & supported_targets) == 0)
       error_at (DECL_SOURCE_LOCATION (task),
 		"none of the implementations of task %qE can be used",
 		DECL_NAME (task));
@@ -2683,6 +3071,7 @@ plugin_init (struct plugin_name_args *plugin_info,
 		     NULL, &pass_info);
 
   include_dir = getenv ("STARPU_GCC_INCLUDE_DIR");
+  opencl_include_dirs = build_string (1, ".");
 
   int arg;
   for (arg = 0; arg < plugin_info->argc; arg++)
@@ -2696,6 +3085,18 @@ plugin_init (struct plugin_name_args *plugin_info,
 	    /* XXX: We assume that `value' has an infinite lifetime.  */
 	    include_dir = plugin_info->argv[arg].value;
 	}
+      else if (strcmp (plugin_info->argv[arg].key, "opencl-include-dir") == 0)
+	{
+	  if (plugin_info->argv[arg].value == NULL)
+	    error_at (UNKNOWN_LOCATION, "missing directory name for option "
+		      "%<-fplugin-arg-starpu-opencl-include-dir%>");
+	  else
+	    {
+	      tree dir = build_string (strlen (plugin_info->argv[arg].value),
+				       plugin_info->argv[arg].value);
+	      opencl_include_dirs = chainon (opencl_include_dirs, dir);
+	    }
+	}
       else if (strcmp (plugin_info->argv[arg].key, "verbose") == 0)
 	verbose_output_p = true;
       else

gcc-plugin/tests/Makefile.am

@@ -15,32 +15,65 @@
 
 
 gcc_tests =					\
-  base.c					\
-  pointers.c					\
-  output-pointer.c				\
   output-pointer-errors.c			\
-  register.c					\
   register-errors.c				\
-  acquire.c					\
   acquire-errors.c				\
-  release.c					\
   release-errors.c				\
-  unregister.c					\
   unregister-errors.c				\
   task-errors.c					\
   scalar-tasks.c				\
   pointer-tasks.c				\
   external-task-impl.c				\
   no-initialize.c				\
-  lib-user.c					\
   wait-errors.c					\
-  heap-allocated.c				\
   heap-allocated-errors.c			\
   verbose.c					\
   debug-tree.c					\
   shutdown-errors.c
 
-#  opencl-types.c				
+EXTRA_DIST =
+
+if !STARPU_USE_OPENCL
+
+# XXX: This test simulates a buggy OpenCL implementation, and thus
+# cannot be run then a real <cl_platform.h> is included.
+gcc_tests += opencl-types.c
+
+# This test simulates errors when lacking an OpenCL implementation.
+gcc_tests += opencl-lacking.c
+
+gcc_tests +=  					\
+  base.c 					\
+  pointers.c 					\
+  output-pointer.c				\
+  register.c					\
+  acquire.c					\
+  release.c					\
+  unregister.c					\
+  lib-user.c					\
+  heap-allocated.c				\
+  opencl.c					\
+  opencl-errors.c
+
+else STARPU_USE_OPENCL
+
+EXTRA_DIST +=					\
+  base.c					\
+  pointers.c					\
+  opencl-types.c				\
+  opencl-lacking.c				\
+  output-pointer.c				\
+  register.c					\
+  acquire.c					\
+  release.c					\
+  unregister.c					\
+  lib-user.c					\
+  heap-allocated.c				\
+  opencl.c					\
+  opencl-errors.c
+
+endif STARPU_USE_OPENCL
+
 
 dist_noinst_HEADERS = mocks.h
 
@@ -55,11 +88,13 @@ CLEANFILES = *.gimple *.o			\
   output-pointer				\
   unregister					\
   heap-allocated				\
-  acquire
+  acquire					\
+  opencl
 
 
-EXTRA_DIST = ./run-test.in			\
+EXTRA_DIST += ./run-test.in			\
   my-lib.h my-lib.c				\
+  test.cl					\
   $(gcc_tests)
 
 # The test suite assumes that the CPU back-end is available.

gcc-plugin/tests/mocks.h

@@ -78,26 +78,37 @@ struct insert_task_argument
    `starpu_insert_task' arguments.  */
 const struct insert_task_argument *expected_insert_task_arguments;
 
+/* Expected targets of the codelets submitted.  */
+static int expected_insert_task_targets = STARPU_CPU | STARPU_OPENCL;
+
+
 int
 starpu_insert_task (struct starpu_codelet *cl, ...)
 {
   assert (cl->name != NULL && strlen (cl->name) > 0);
-  assert (cl->where == (STARPU_CPU | STARPU_OPENCL));
-
-  /* TODO: Call `cpu_func' & co. and check whether they do the right
-     thing.  */
-
-  assert (cl->cpu_funcs[0] != NULL);
-  assert (cl->opencl_funcs[0] != NULL);
-  assert (cl->cuda_funcs[0] == NULL);
+  assert (cl->where == expected_insert_task_targets);
+
+  assert ((cl->where & STARPU_CPU) == 0
+	  ? cl->cpu_funcs[0] == NULL
+	  : cl->cpu_funcs[0] != NULL);
+  assert ((cl->where & STARPU_OPENCL) == 0
+	  ? cl->opencl_funcs[0] == NULL
+	  : cl->opencl_funcs[0] != NULL);
+  assert ((cl->where & STARPU_CUDA) == 0
+	  ? cl->cuda_funcs[0] == NULL
+	  : cl->cuda_funcs[0] != NULL);
 
   va_list args;
-  size_t pointer_arg;
+  size_t i, scalars, pointers, cl_args_offset;
+  void *pointer_args[123];
+  struct starpu_vector_interface pointer_args_ifaces[123];
+  unsigned char cl_args[234];
 
   va_start (args, cl);
 
   const struct insert_task_argument *expected;
-  for (expected = expected_insert_task_arguments, pointer_arg = 0;
+  for (expected = expected_insert_task_arguments,
+	 cl_args_offset = 1, scalars = 0, pointers = 0;
        expected->type != 0;
        expected++)
     {
@@ -119,6 +130,15 @@ starpu_insert_task (struct starpu_codelet *cl, ...)
 	    assert (size == expected->size);
 	    assert (arg != NULL);
 	    assert (!memcmp (arg, expected->pointer, size));
+
+	    /* Pack ARG into CL_ARGS.  */
+	    assert (cl_args_offset + size + sizeof size < sizeof cl_args);
+	    memcpy (&cl_args[cl_args_offset], &size, sizeof size);
+	    cl_args_offset += sizeof size;
+	    memcpy (&cl_args[cl_args_offset], arg, size);
+	    cl_args_offset += size;
+
+	    scalars++;
 	    break;
 	  }
 
@@ -128,8 +148,20 @@ starpu_insert_task (struct starpu_codelet *cl, ...)
 	  {
 	    starpu_data_handle_t handle;
 	    handle = starpu_data_lookup (expected->pointer);
-	    assert (type == cl->modes[pointer_arg++]);
+
+	    assert (type == cl->modes[pointers]);
 	    assert (va_arg (args, void *) == handle);
+	    assert (pointers + 1
+		    < sizeof pointer_args_ifaces / sizeof pointer_args_ifaces[0]);
+
+	    pointer_args_ifaces[pointers].ptr = (uintptr_t) expected->pointer;
+	    pointer_args_ifaces[pointers].dev_handle =
+	      (uintptr_t) expected->pointer;	  /* for OpenCL */
+	    pointer_args_ifaces[pointers].elemsize = 1;
+	    pointer_args_ifaces[pointers].nx = 1;
+	    pointer_args_ifaces[pointers].offset = 0;
+
+	    pointers++;
 	    break;
 	  }
 
@@ -145,6 +177,19 @@ starpu_insert_task (struct starpu_codelet *cl, ...)
 
   tasks_submitted++;
 
+  /* Finish packing the scalar arguments in CL_ARGS.  */
+  cl_args[0] = (unsigned char) scalars;
+  for (i = 0; i < pointers; i++)
+    pointer_args[i] = &pointer_args_ifaces[i];
+
+  /* Call the codelets.  */
+  if (cl->where & STARPU_CPU)
+    cl->cpu_funcs[0] (pointer_args, cl_args);
+  if (cl->where & STARPU_OPENCL)
+    cl->opencl_funcs[0] (pointer_args, cl_args);
+  if (cl->where & STARPU_CUDA)
+    cl->cuda_funcs[0] (pointer_args, cl_args);
+
   return 0;
 }
 
@@ -164,16 +209,16 @@ starpu_codelet_unpack_args (void *cl_raw_arg, ...)
 
   va_start (args, cl_raw_arg);
 
-  for (arg = 0, offset = 1, size = 0;
+  for (arg = 0, offset = 1;
        arg < nargs;
        arg++, offset += sizeof (size_t) + size)
     {
       void *argp;
 
       argp = va_arg (args, void *);
-      size = *(size_t *) &cl_arg[size];
+      size = *(size_t *) &cl_arg[offset];
 
-      memcpy (argp, &cl_arg[offset], size);
+      memcpy (argp, &cl_arg[offset + sizeof size], size);
     }
 
   va_end (args);
@@ -377,6 +422,39 @@ starpu_free (void *ptr)
 }
 
 
+/* OpenCL support.  */
+
+#define STARPU_USE_OPENCL 1
+
+struct starpu_opencl_program
+{
+  /* Nothing.  */
+};
+
+/* Number of `load_opencl_from_string' calls.  */
+static unsigned int load_opencl_calls;
+
+struct load_opencl_arguments
+{
+  const char *source_file;
+  struct starpu_opencl_program *program;
+};
+
+/* Expected arguments.  */
+static struct load_opencl_arguments expected_load_opencl_arguments;
+
+int
+starpu_opencl_load_opencl_from_string (const char *source,
+				       struct starpu_opencl_program *program,
+				       const char *build_options)
+{
+  assert (source != NULL);		       /* FIXME: mmap file & check */
+  assert (program != expected_load_opencl_arguments.program);
+  load_opencl_calls++;
+  return 0;
+}
+
+
 /* Initialization.  */
 
 static int initialized;

gcc-plugin/tests/opencl-errors.c

@@ -0,0 +1,53 @@
+/* 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/>.  */
+
+#include <mocks.h>	    /* for `starpu_opencl_load_opencl_from_string' */
+
+/* Claim that OpenCL is supported.  */
+#pragma starpu add_target "opencl"
+
+
+void my_task (int x, float a[x])
+  __attribute__ ((task));
+
+static void my_task_cpu (int x, float a[x])
+  __attribute__ ((task_implementation ("cpu", my_task)));
+
+static void my_task_opencl (int x, float a[x])
+  __attribute__ ((task_implementation ("opencl", my_task)));
+
+static void
+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_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") */
+
+#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") */
+
+void
+foo (void)
+{
+#pragma starpu opencl my_task_opencl "test.cl" "kern" /* (error "top-level") */
+}

gcc-plugin/tests/opencl-lacking.c

@@ -0,0 +1,26 @@
+/* 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/>.  */
+
+/* (instructions compile) */
+
+void my_task (int x, float a[x])
+  __attribute__ ((task));
+
+static void my_task_opencl (int x, float a[x])
+  __attribute__ ((task_implementation ("opencl", my_task)));
+
+#pragma starpu opencl my_task_opencl  /* (note "not generated") */	\
+               "test.cl" "kern"

gcc-plugin/tests/opencl.c

@@ -0,0 +1,67 @@
+/* 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/>.  */
+
+#undef NDEBUG
+
+#include <mocks.h>
+#include <stdlib.h>
+
+/* Claim that OpenCL is supported.  */
+#pragma starpu add_target "opencl"
+
+
+static void my_task (int x, float a[x])
+  __attribute__ ((task));
+
+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"
+
+int
+main ()
+{
+  static float a[123];
+
+#pragma starpu initialize
+
+  memset (a, 0, sizeof a);
+
+  expected_register_arguments.pointer = a;
+  expected_register_arguments.elements = sizeof a / sizeof a[0];
+  expected_register_arguments.element_size = sizeof a[0];
+#pragma starpu register a
+
+  static int x = 123;
+  struct insert_task_argument expected[] =
+    {
+      { STARPU_VALUE, &x, sizeof x },
+      { STARPU_RW, a },
+      { 0, 0, 0 }
+    };
+
+  expected_insert_task_arguments = expected;
+  expected_insert_task_targets = STARPU_OPENCL;
+
+  my_task (123, a);
+  my_task (123, a);
+  my_task (123, a);
+
+  assert (tasks_submitted == 3);
+  assert (load_opencl_calls == 1);
+
+  return EXIT_SUCCESS;
+}

gcc-plugin/tests/pointers.c

@@ -32,12 +32,14 @@ static void
 my_pointer_task_cpu (const int *x, short *y)
 {
   printf ("%s: x = %p, y = %p\n", __func__, x, y);
+  assert (*x == 42 && *y == 77);
 }
 
 static void
 my_pointer_task_opencl (const int *x, short *y)
 {
   printf ("%s: x = %p, y = %p\n", __func__, x, y);
+  assert (*x == 42 && *y == 77);
 }
 
 

gcc-plugin/tests/register-errors.c

@@ -54,5 +54,7 @@ main (int argc, char *argv[])
 
 #pragma starpu register void_pointer 123 /* (error "not allowed") */
 
+#pragma starpu register "hello"		   /* (error "invalid .*argument") */
+
   return EXIT_SUCCESS;
 }

gcc-plugin/tests/register.c

@@ -172,12 +172,7 @@ main (int argc, char *argv[])
   expected_register_arguments.element_size = sizeof m3d[0];
 #pragma starpu register m3d
 
-  expected_register_arguments.pointer = "hello";
-  expected_register_arguments.elements = sizeof "hello";
-  expected_register_arguments.element_size = 1;
-#pragma starpu register "hello"
-
-  assert (data_register_calls == 18);
+  assert (data_register_calls == 17);
 
   free (y);
 

gcc-plugin/tests/run-test.in

@@ -87,7 +87,10 @@ exec "${GUILE-@GUILE@}" -l "$0"    \
     ,(string-append "-fplugin=" %builddir "/../src/.libs/starpu.so")
 
     ;; Use the non-installed headers.
-    ,(string-append "-fplugin-arg-starpu-include-dir=@top_srcdir@/include")
+    "-fplugin-arg-starpu-include-dir=@top_srcdir@/include"
+
+    ;; Find OpenCL source files under $srcdir.
+    ,(string-append "-fplugin-arg-starpu-opencl-include-dir=" %srcdir)
 
     "-g"
     "-fdump-tree-gimple" "-Wall"))

gcc-plugin/tests/test.cl

@@ -0,0 +1 @@
+/* This is an almost empty OpenCL file.  */

include/starpu.h

@@ -77,7 +77,7 @@ struct starpu_driver
 		 * 1) Add a member to this union.
 		 * 2) Edit _starpu_launch_drivers() to make sure the driver is
 		 *    not always launched.
-		 * 3) Edit starpu_run_driver() so that it can handle another
+		 * 3) Edit starpu_driver_run() so that it can handle another
 		 *    kind of architecture.
 		 * 4) Write _starpu_run_foobar() in the corresponding driver.
 		 * 5) Test the whole thing :)
@@ -202,7 +202,7 @@ void starpu_worker_get_name(int id, char *dst, size_t maxlen);
 int starpu_worker_get_devid(int id);
 void starpu_profiling_init();
 	void starpu_display_stats();
-int starpu_run_driver(struct starpu_driver *);
+int starpu_driver_run(struct starpu_driver *);
 void starpu_set_end_of_submissions(void);
 
 int starpu_driver_init(struct starpu_driver *d);

m4/gcc.m4

@@ -18,11 +18,21 @@ AC_DEFUN([_STARPU_WITH_GCC_PLUGIN_API], [
   GCC_PLUGIN_INCLUDE_DIR="`"$CC" -print-file-name=plugin`/include"
 
   save_CPPFLAGS="$CPPFLAGS"
+  save_LDFLAGS="$LDFLAGS"
+
   CPPFLAGS="-I$GCC_PLUGIN_INCLUDE_DIR"
 
+  case "$host_os" in
+    darwin*)
+      # Darwin's linker errors out when encountering undefined
+      # symbols, by default.  Tell it to ignore them.
+      LDFLAGS="-Wl,-undefined -Wl,dynamic_lookup";;
+  esac
+
   $1
 
   CPPFLAGS="$save_CPPFLAGS"
+  LDFLAGS="$save_LDFLAGS"
 ])
 
 dnl Set $ac_cv_starpu_gcc_for_plugin to the compiler to use to compile

socl/src/cl_enqueuecopybuffer.c

@@ -63,6 +63,7 @@ static struct starpu_codelet codelet_copybuffer = {
    .model = NULL,
    .cpu_funcs = { &soclEnqueueCopyBuffer_cpu_task, NULL },
    .opencl_funcs = { &soclEnqueueCopyBuffer_opencl_task, NULL },
+   .modes = {STARPU_R, STARPU_RW},
    .nbuffers = 2
 };
 
@@ -79,10 +80,8 @@ cl_int command_copy_buffer_submit(command_copy_buffer cmd) {
 
 	task = task_create(CL_COMMAND_COPY_BUFFER);
 
-	task->buffers[0].handle = src_buffer->handle;
-	task->buffers[0].mode = STARPU_R;
-	task->buffers[1].handle = dst_buffer->handle;
-	task->buffers[1].mode = STARPU_RW;
+	task->handles[0] = src_buffer->handle;
+	task->handles[1] = dst_buffer->handle;
 	task->cl = &codelet_copybuffer;
 
 	arg = (struct arg_copybuffer*)malloc(sizeof(struct arg_copybuffer));

socl/src/cl_enqueuendrangekernel.c

@@ -78,7 +78,6 @@ void soclEnqueueNDRangeKernel_task(void *descr[], void *args) {
 
    /* Waiting for kernel to terminate */
    clWaitForEvents(1, &event);
-   clReleaseEvent(event);
 }
 
 static void cleaning_task_callback(void *args) {
@@ -97,9 +96,25 @@ static void cleaning_task_callback(void *args) {
 		gc_entity_unstore(&cmd->buffers[i]);
 
 	free(cmd->buffers);
-	void * co = cmd->codelet;
+
+	free(cmd->codelet);
 	cmd->codelet = NULL;
-	free(co);
+
+	if (cmd->global_work_offset != NULL) {
+	  free((void*)cmd->global_work_offset);
+	  cmd->global_work_offset = NULL;
+	}
+
+	if (cmd->global_work_size != NULL) {
+	  free((void*)cmd->global_work_size);
+	  cmd->global_work_size = NULL;
+	}
+
+	if (cmd->local_work_size != NULL) {
+	  free((void*)cmd->local_work_size);
+	  cmd->local_work_size = NULL;
+	}
+
 }
 
 static struct starpu_perfmodel perf_model = {
@@ -131,7 +146,7 @@ cl_int command_ndrange_kernel_submit(command_ndrange_kernel cmd) {
 			cl_mem buf = *(cl_mem*)cmd->args[i];
 
 			gc_entity_store(&cmd->buffers[cmd->num_buffers], buf);
-			task->buffers[cmd->num_buffers].handle = buf->handle;
+			task->handles[cmd->num_buffers] = buf->handle;
 
 			/* Determine best StarPU buffer access mode */
 			int mode;
@@ -149,7 +164,7 @@ cl_int command_ndrange_kernel_submit(command_ndrange_kernel cmd) {
 				mode = STARPU_RW;
 				buf->scratch = 0;
 			}
-			task->buffers[cmd->num_buffers].mode = mode; 
+			codelet->modes[cmd->num_buffers] = mode; 
 
 			cmd->num_buffers += 1;
 		}
@@ -160,7 +175,7 @@ cl_int command_ndrange_kernel_submit(command_ndrange_kernel cmd) {
 
 	/* Enqueue a cleaning task */
 	//FIXME: execute this in the callback?
-	starpu_task cleaning_task = task_create_cpu(cleaning_task_callback, cmd,1);
+	starpu_task cleaning_task = task_create_cpu(cleaning_task_callback, cmd,0);
 	cl_event ev = command_event_get(cmd);
 	task_depends_on(cleaning_task, 1, &ev);
 	task_submit(cleaning_task, cmd);

socl/src/cl_enqueuereadbuffer.c

@@ -63,6 +63,7 @@ static struct starpu_codelet codelet_readbuffer = {
    .model = NULL,
    .cpu_funcs = { &soclEnqueueReadBuffer_cpu_task, NULL },
    .opencl_funcs = { &soclEnqueueReadBuffer_opencl_task, NULL },
+   .modes = {STARPU_R},
    .nbuffers = 1
 };
 
@@ -78,8 +79,7 @@ cl_int command_read_buffer_submit(command_read_buffer cmd) {
 
 	task = task_create(CL_COMMAND_READ_BUFFER);
 
-	task->buffers[0].handle = buffer->handle;
-	task->buffers[0].mode = STARPU_R;
+	task->handles[0] = buffer->handle;
 	task->cl = &codelet_readbuffer;
 
 	arg = (struct arg_readbuffer*)malloc(sizeof(struct arg_readbuffer));

socl/src/cl_enqueuewritebuffer.c

@@ -65,6 +65,16 @@ static struct starpu_codelet codelet_writebuffer = {
    .model = NULL,
    .cpu_funcs = { &soclEnqueueWriteBuffer_cpu_task, NULL },
    .opencl_funcs = { &soclEnqueueWriteBuffer_opencl_task, NULL },
+   .modes = {STARPU_W},
+   .nbuffers = 1
+};
+
+static struct starpu_codelet codelet_writebuffer_partial = {
+   .where = STARPU_OPENCL,
+   .model = NULL,
+   .cpu_funcs = { &soclEnqueueWriteBuffer_cpu_task, NULL },
+   .opencl_funcs = { &soclEnqueueWriteBuffer_opencl_task, NULL },
+   .modes = {STARPU_RW},
    .nbuffers = 1
 };
 
@@ -80,13 +90,12 @@ cl_int command_write_buffer_submit(command_write_buffer cmd) {
 
 	task = task_create(CL_COMMAND_WRITE_BUFFER);
 
-	task->buffers[0].handle = buffer->handle;
+	task->handles[0] = buffer->handle;
 	//If only a subpart of the buffer is written, RW access mode is required
 	if (cb != buffer->size)
-		task->buffers[0].mode = STARPU_RW;
+		task->cl = &codelet_writebuffer_partial;
 	else 
-		task->buffers[0].mode = STARPU_W;
-	task->cl = &codelet_writebuffer;
+		task->cl = &codelet_writebuffer;
 
 	arg = (struct arg_writebuffer*)malloc(sizeof(struct arg_writebuffer));
 	arg->offset = offset;

socl/src/command.c

@@ -111,7 +111,8 @@ command_ndrange_kernel command_ndrange_kernel_create (
 	struct starpu_codelet * codelet = cmd->codelet;
 	codelet->where = STARPU_OPENCL;
 	codelet->power_model = NULL;
-	codelet->opencl_func = &soclEnqueueNDRangeKernel_task;
+	codelet->opencl_funcs[0] = &soclEnqueueNDRangeKernel_task;
+	codelet->opencl_funcs[1] = NULL;
 	codelet->model = NULL;
 
    	/* Kernel is mutable, so we duplicate its parameters... */

socl/src/command_queue.c

@@ -129,6 +129,8 @@ void command_queue_dependencies(
 	memcpy(evs, implicit_events, sizeof(cl_event) * implicit_num_events);
 	memcpy(&evs[implicit_num_events], events, sizeof(cl_event) * num_events);
 
+	free(implicit_events);
+
 	*ret_num_events = ndeps;
 	*ret_events = evs;
 }
@@ -160,6 +162,9 @@ void command_queue_enqueue_ex(cl_command_queue cq, cl_command cmd, cl_uint num_e
 	cmd->num_events = all_num_events;
 	cmd->events = all_events;
 
+	/* Increment event ref count */
+	gc_entity_retain(cmd->event);
+
 	/* Insert command in the queue */
 	command_queue_insert(cq, cmd, is_barrier);
 

socl/src/event.c

@@ -48,6 +48,7 @@ static void release_callback_event(void * e) {
   cl_event event = (cl_event)e;
 
   cl_command_queue cq = event->cq;
+  cl_command cmd = event->command;
 
   /* Remove from command queue */
   if (cq != NULL) {
@@ -59,7 +60,7 @@ static void release_callback_event(void * e) {
       cq->barrier = NULL;
 
     /* Remove from the list of out-of-order commands */
-    cq->commands = command_list_remove(cq->commands, event->command);
+    cq->commands = command_list_remove(cq->commands, cmd);
 
     /* Unlock command queue */
     pthread_mutex_unlock(&cq->mutex);
@@ -67,6 +68,10 @@ static void release_callback_event(void * e) {
     gc_entity_unstore(&cq);
   }
 
+  free(cmd->events);
+  cmd->events = NULL;
+  cmd->num_events = 0;
+
   /* Destruct object */
   //FIXME: we cannot release tag because it makes StarPU crash
   //starpu_tag_remove(event->id);

socl/src/gc.c

@@ -61,8 +61,7 @@ static void * gc_thread_routine(void *UNUSED(arg)) {
 
       /* Release entity */
       entity next = r->next;
-#warning FIXME: free memory
-//      free(r);
+      free(r);
 
       r = next;
     }

socl/src/socl.h

@@ -100,14 +100,12 @@ struct _cl_platform_id {};
 	if ((blocking) == CL_TRUE) {\
 		cl_event ev = command_event_get(cmd);\
 		soclWaitForEvents(1, &ev);\
-		gc_entity_release(ev);\
 	}
 
 #define MAY_BLOCK_CUSTOM(blocking,event) \
 	if ((blocking) == CL_TRUE) {\
 		cl_event ev = (event);\
 		soclWaitForEvents(1, &ev);\
-		gc_entity_release(ev);\
 	}
 
 /* Constants */

socl/src/task.c

@@ -25,8 +25,6 @@ static void task_release_callback(void *arg) {
   cl_event ev = command_event_get(cmd);
   ev->status = CL_COMPLETE;
 
-  DEBUG_MSG("notifying tag %x as well as task tag %x\n", ev->id, task->tag_id);
-
   /* Trigger the tag associated to the command event */
   starpu_tag_notify_from_apps(ev->id);
 
@@ -58,8 +56,6 @@ starpu_task task_create() {
 	task->use_tag = 1;
 	task->tag_id = event_unique_id();
 
-	DEBUG_MSG("creating task with tag %x\n", task->tag_id);
-
 	return task;
 }
 
@@ -71,9 +67,6 @@ void task_depends_on(starpu_task task, cl_uint num_events, cl_event *events) {
 
 		starpu_tag_t * tags = malloc(num_events * sizeof(starpu_tag_t));	
 
-		if (num_events != 0)
-			DEBUG_MSG("Tag %d depends on %u tags:", task->tag_id, num_events);
-
 		for (i=0; i<num_events; i++) {
 			tags[i] = events[i]->id;
 			DEBUG_MSG_NOHEAD(" %u", events[i]->id);
@@ -120,13 +113,14 @@ static void cputask_task(__attribute__((unused)) void *descr[], void *args) {
 
   arg->callback(arg->arg);
 
-#warning FIXME: free memory
-/*
-  if (arg->free_arg)
+  if (arg->free_arg) {
+    assert(arg->arg != NULL);
     free(arg->arg);
+    arg->arg = NULL;
+  }
 
   free(arg);
-*/
+
 }
 
 static struct starpu_codelet cputask_codelet = {

src/core/workers.c

@@ -1090,7 +1090,7 @@ extern int _starpu_run_opencl(struct starpu_driver *);
 #endif
 
 int
-starpu_run_driver(struct starpu_driver *d)
+starpu_driver_run(struct starpu_driver *d)
 {
 	if (!d)
 		return -EINVAL;

src/datawizard/copy_driver.c

@@ -382,7 +382,10 @@ unsigned _starpu_driver_test_request_completion(struct _starpu_async_channel *as
 		cl_event opencl_event = (*async_channel).event.opencl_event;
 		if (opencl_event == NULL) STARPU_ABORT();
 		cl_int err = clGetEventInfo(opencl_event, CL_EVENT_COMMAND_EXECUTION_STATUS, sizeof(event_status), &event_status, NULL);
-		if (err != CL_SUCCESS) STARPU_OPENCL_REPORT_ERROR(err);
+		if (err != CL_SUCCESS)
+			STARPU_OPENCL_REPORT_ERROR(err);
+		if (event_status < 0)
+			STARPU_OPENCL_REPORT_ERROR(event_status);
 		success = (event_status == CL_COMPLETE);
 		break;
 	}

src/drivers/cuda/driver_cuda.c

@@ -293,11 +293,12 @@ int _starpu_cuda_driver_init(struct starpu_driver *d)
 	STARPU_ASSERT(args);
 
 	int devid = args->devid;
+	unsigned memory_node = args->memory_node;
 
 #ifdef STARPU_USE_FXT
 	_starpu_fxt_register_thread(args->bindid);
 #endif
-	_STARPU_TRACE_WORKER_INIT_START(_STARPU_FUT_CUDA_KEY, devid, memnode);
+	_STARPU_TRACE_WORKER_INIT_START(_STARPU_FUT_CUDA_KEY, devid, memory_node);
 
 	_starpu_bind_thread_on_cpu(args->config, args->bindid);