This is the mail archive of the gcc-patches@gcc.gnu.org mailing list for the GCC project.


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]
Other format: [Raw text]

Re: [PATCH,openacc] check for compatible loop parallelism with acc routine calls


On 06/17/2016 07:42 AM, Jakub Jelinek wrote:
> On Wed, Jun 15, 2016 at 08:12:15PM -0700, Cesar Philippidis wrote:
>> The second set of changes involves teaching the gimplifier to error when
>> it detects a function call to an non-acc routines inside an OpenACC
>> offloaded region. Actually, I relaxed non-acc routines by excluding
>> calls to builtin functions, including those prefixed with _gfortran_.
>> Nvptx does have a newlib c library, and it also has a subset of
>> libgfortran. Still, this solution is probably not optimal.
> 
> I don't really like that, hardcoding prefixes or whatever is available
> (you have quite some subset of libc, libm etc. available too) in the
> compiler looks very hackish.  What is wrong with complaining during
> linking of the offloaded code?

Wouldn't the error get reported multiple times then, i.e. once per
target? Then again, maybe this error could have been restrained to the
host compiler.

Anyway, this patch now reduces that error to a warning. Furthermore,
that warning is being thrown in lower_omp_1 instead of
gimplify_call_expr because the latter is called multiple times and that
causes duplicate warnings. The only bit of fallout I had with this
change was with the fortran FE's usage of BUILT_IN_EXPECT in
gfc_{un}likely. Since these are generated implicitly by the FE, I just
added an oacc_function attribute to those calls when flag_openacc is set.

>> Next, I had to modify the openacc header files in libgomp to mark
>> acc_on_device as an acc routine. Unfortunately, this meant that I had to
>> build the opeancc.mod module for gfortran with -fopenacc. But doing
>> that, caused caused gcc to stream offloaded code to the openacc.o object
>> file. So, I've updated the behavior of flag_generate_offload such that
>> minus one indicates that the user specified -foffload=disable, and that
>> will prevent gcc from streaming offloaded lto code. The alternative was
>> to hack libtool to build libgomp with -foffload=disable.
> 
> This also looks wrong.  I'd say the right thing is when loading modules
> that have OpenACC bits set in it (and also OpenMP bits, I admit I haven't
> handled this well) into CU with the corresponding flags unset (-fopenacc,
> -fopenmp, -fopenmp-simd here, depending on which bit it is), then
> IMHO the module loading code should just ignore it, pretend it wasn't there.
> Similarly e.g. to how lto1 with -g0 should ignore debug statements that
> could be in the LTO inputs.

This required two changes. First, I had to teach lto-cgraph.c how to
report an error rather then fail an assert when partitions are missing
decls. Second, I taught the lto wrapper how to stream offloaded code on
the absence of -fopen*. The only kink with this approach is that I had
to build libgomp/openacc.f90 with -frandom-seed=1 to prevent lto related
bootstrap failures.

By the way, Thomas, I've added

 #pragma acc routine(__builtin_acc_on_device) seq

to openacc.h. Is this OK, or should I just modify the various
libgomp.oacc-c-c++-common/loop* tests to use that pragma directly? Or
another option is to have the compiler add that attribute directly. I
don't think we're really expecting the end user to use
__builtin_acc_on_device directly since this is a gcc-ism.

Cesar
2016-06-23  Cesar Philippidis  <cesar@codesourcery.com>

	gcc/
	* lto-cgraph.c (input_overwrite_node): Error on missing symbols.
	(input_varpool_node): Likewise.
	* lto-wrapper.c (compile_images_for_offload_targets): Don't stream
	offloaded images without -fopenacc, -fopenmp or -fopenmp-simd.
	(run_gcc): Set flag_openacc, flag_openmp, and flag_openmp_simd.
	* omp-low.c (lower_omp_1): Emit a warning when calling a function
	that doesn't have an oacc_function attribute from an OpenACC offloaded
	region.
	(oacc_loop_fixed_partitions): Consider SEQ loops when checking
	parallelism.

	gcc/fortran/
	* gfortran.h (enum oacc_function): New enum.
	(oacc_function_types): Declare.
	(symbol_attribute): Add oacc_function field.
	(gfc_intrinsic_sym): Likewise.
	(add_omp_offloading_attributes): Declare.
	* intrinsic.c (add_sym): Initialize oacc_fuction to zero.
	(gfc_intrinsic_sub_interface): Set attr.oacc_function as to
	OACC_FUNCTION_SEQ in the resolved symbol when appropriate.
	* module.c (oacc_function): New DECL_MIO_NAME.
	(mio_symbol_attribute): Set attr->oacc_function.
	* openmp.c (gfc_oacc_routine_dims): Change return type to oacc_function.
	(gfc_match_oacc_routine): Permit named 'acc routine' directives on
	intrinsic procedures.  Update call to gfc_oacc_routine_dims.
	* symbol.c (oacc_function_types): Define.
	* trans-decl.c (add_omp_offloading_attributes): New function.
	(add_attributes_to_decl): Use it.
	* trans.c (gfc_unlikely): Mark calls BUILT_IN_EXPECT as 'acc routines'
	with flag_openacc is set.
	(gfc_likely): Likewise.

	gcc/testsuite/
	* c-c++-common/goacc/kernels-1.c: Add warnings to calls to
	__builtin_abort.
	* c-c++-common/goacc/parallel-1.c: Likewise.
	* c-c++-common/goacc/routine-3.c: Add coverage for acc seq loops.
	* c-c++-common/goacc/routine-6.c: New test.
	* gfortran.dg/goacc/fixed-1.f: Mark abort as an 'acc routine'.
	* gfortran.dg/goacc/routine-7.f90: New test.
	* gfortran.dg/goacc/routine-8.f90: New test.

	libgomp/
	* Makefile.am (openacc.lo): New rule.
	(openacc.mod): Build with -fopenacc -frandom-seed=1.
	* Makefile.in: Regenerate.
	* openacc.f90 (acc_on_device_h): Mark as 'acc routine seq'.
	(acc_on_device_l): Likewise.
	* openacc.h (acc_on_device): Mark as 'acc routine seq'.
	(__builtin_acc_on_device): New declaration. Mark as 'acc routine seq'.
	* openacc_lib.h (acc_on_device_h): Mark as 'acc routine seq'.
	* testsuite/libgomp.oacc-c-c++-common/abort-1.c: Apply 'acc routine
	seq' on abort.
	* testsuite/libgomp.oacc-c-c++-common/abort-1.c: Add pragma 'acc
	routine(abort) seq'.
	* testsuite/libgomp.oacc-c-c++-common/abort-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/abort-3.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/abort-4.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/abort-5.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/acc_on_device-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/data-clauses-kernels-ipa-pta.c:
	Add -fno-exceptions to dg-additional-options.
	* testsuite/libgomp.oacc-c-c++-common/data-clauses-kernels.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/data-clauses-parallel-ipa-pta.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/data-clauses-parallel.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/data-clauses.h: Add pragma 'acc
	routine(__builtin_abort) seq'.
	* testsuite/libgomp.oacc-c-c++-common/if-1.c: Add pragma 'acc
	routine(abort) seq'.
	* testsuite/libgomp.oacc-c-c++-common/loop-g-1.c: Include openacc.h.
	pass acc_device_nvidia to __builtin_acc_on_device.
	* testsuite/libgomp.oacc-c-c++-common/loop-g-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-red-gwv-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-red-v-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-red-wv-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-v-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-w-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/routine-4.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/routine-g-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/routine-v-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/routine-w-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/vector-type-1.c: Add pragma
	'acc routine(__builtin_abort) seq'.
	* testsuite/libgomp.oacc-fortran/abort-1.f90: Add directive 'acc
	routine(abort) seq'.
	* testsuite/libgomp.oacc-fortran/abort-2.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f: Likewise.
	* testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f: Likewise.
	* testsuite/libgomp.oacc-fortran/nested-function-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/routine-7.f90: Update test to be
	thread safe.

diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index 0bb71cb..bf46931 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -303,6 +303,15 @@ enum save_state
 { SAVE_NONE = 0, SAVE_EXPLICIT, SAVE_IMPLICIT
 };
 
+/* Flags to keep track of ACC routine states.  */
+enum oacc_function
+{ OACC_FUNCTION_NONE = 0,
+  OACC_FUNCTION_SEQ,
+  OACC_FUNCTION_GANG,
+  OACC_FUNCTION_WORKER,
+  OACC_FUNCTION_VECTOR
+};
+
 /* Strings for all symbol attributes.  We use these for dumping the
    parse tree, in error messages, and also when reading and writing
    modules.  In symbol.c.  */
@@ -312,6 +321,7 @@ extern const mstring intents[];
 extern const mstring access_types[];
 extern const mstring ifsrc_types[];
 extern const mstring save_status[];
+extern const mstring oacc_function_types[];
 
 /* Enumeration of all the generic intrinsic functions.  Used by the
    backend for identification of a function.  */
@@ -862,7 +872,7 @@ typedef struct
   unsigned oacc_declare_link:1;
 
   /* This is an OpenACC acclerator function at level N - 1  */
-  unsigned oacc_function:3;
+  ENUM_BITFIELD (oacc_function) oacc_function:3;
 
   /* Attributes set by compiler extensions (!GCC$ ATTRIBUTES).  */
   unsigned ext_attr:EXT_ATTR_NUM;
@@ -1956,7 +1966,7 @@ typedef struct gfc_intrinsic_sym
   gfc_typespec ts;
   unsigned elemental:1, inquiry:1, transformational:1, pure:1,
     generic:1, specific:1, actual_ok:1, noreturn:1, conversion:1,
-    from_module:1, vararg:1;
+    from_module:1, vararg:1, oacc_function:1;
 
   int standard;
 
@@ -3299,5 +3309,8 @@ bool gfc_is_reallocatable_lhs (gfc_expr *);
 /* trans-decl.c */
 
 void finish_oacc_declare (gfc_namespace *, gfc_symbol *, bool);
+tree add_omp_offloading_attributes (unsigned omp_declare_target,
+				    enum oacc_function, tree list);
+
 
 #endif /* GCC_GFORTRAN_H  */
diff --git a/gcc/fortran/intrinsic.c b/gcc/fortran/intrinsic.c
index 1d7503d..7b8935b 100644
--- a/gcc/fortran/intrinsic.c
+++ b/gcc/fortran/intrinsic.c
@@ -354,6 +354,7 @@ add_sym (const char *name, gfc_isym_id id, enum klass cl, int actual_ok, bt type
       next_sym->generic = 0;
       next_sym->conversion = 0;
       next_sym->id = id;
+      next_sym->oacc_function = 0;
       break;
 
     default:
@@ -4583,6 +4584,8 @@ gfc_intrinsic_sub_interface (gfc_code *c, int error_flag)
     {
       c->resolved_sym = gfc_get_intrinsic_sub_symbol (isym->lib_name);
       c->resolved_sym->attr.elemental = isym->elemental;
+      if (isym->oacc_function)
+	c->resolved_sym->attr.oacc_function = OACC_FUNCTION_SEQ;
     }
 
   if (gfc_do_concurrent_flag && !isym->pure)
diff --git a/gcc/fortran/module.c b/gcc/fortran/module.c
index 6d3860e..e3ed2a0 100644
--- a/gcc/fortran/module.c
+++ b/gcc/fortran/module.c
@@ -2095,6 +2095,7 @@ DECL_MIO_NAME (procedure_type)
 DECL_MIO_NAME (ref_type)
 DECL_MIO_NAME (sym_flavor)
 DECL_MIO_NAME (sym_intent)
+DECL_MIO_NAME (oacc_function)
 #undef DECL_MIO_NAME
 
 /* Symbol attributes are stored in list with the first three elements
@@ -2116,6 +2117,8 @@ mio_symbol_attribute (symbol_attribute *attr)
   attr->proc = MIO_NAME (procedure_type) (attr->proc, procedures);
   attr->if_source = MIO_NAME (ifsrc) (attr->if_source, ifsrc_types);
   attr->save = MIO_NAME (save_state) (attr->save, save_status);
+  attr->oacc_function = MIO_NAME (oacc_function) (attr->oacc_function,
+						  oacc_function_types);
 
   ext_attr = attr->ext_attr;
   mio_integer ((int *) &ext_attr);
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index f514866..a8446fe 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -1664,21 +1664,31 @@ gfc_match_oacc_cache (void)
 
 /* Determine the loop level for a routine.   */
 
-static int
+static oacc_function
 gfc_oacc_routine_dims (gfc_omp_clauses *clauses)
 {
   int level = -1;
+  oacc_function ret = OACC_FUNCTION_SEQ;
 
   if (clauses)
     {
       unsigned mask = 0;
 
       if (clauses->gang)
-	level = GOMP_DIM_GANG, mask |= GOMP_DIM_MASK (level);
+	{
+	  level = GOMP_DIM_GANG, mask |= GOMP_DIM_MASK (level);
+	  ret = OACC_FUNCTION_GANG;
+	}
       if (clauses->worker)
-	level = GOMP_DIM_WORKER, mask |= GOMP_DIM_MASK (level);
+	{
+	  level = GOMP_DIM_WORKER, mask |= GOMP_DIM_MASK (level);
+	  ret = OACC_FUNCTION_WORKER;
+	}
       if (clauses->vector)
-	level = GOMP_DIM_VECTOR, mask |= GOMP_DIM_MASK (level);
+	{
+	  level = GOMP_DIM_VECTOR, mask |= GOMP_DIM_MASK (level);
+	  ret = OACC_FUNCTION_VECTOR;
+	}
       if (clauses->seq)
 	level = GOMP_DIM_MAX, mask |= GOMP_DIM_MASK (level);
 
@@ -1689,7 +1699,7 @@ gfc_oacc_routine_dims (gfc_omp_clauses *clauses)
   if (level < 0)
     level = GOMP_DIM_MAX;
 
-  return level;
+  return ret;
 }
 
 match
@@ -1700,6 +1710,7 @@ gfc_match_oacc_routine (void)
   match m;
   gfc_omp_clauses *c = NULL;
   gfc_oacc_routine_name *n = NULL;
+  gfc_intrinsic_sym *isym = NULL;
 
   old_loc = gfc_current_locus;
 
@@ -1717,12 +1728,16 @@ gfc_match_oacc_routine (void)
   if (m == MATCH_YES)
     {
       char buffer[GFC_MAX_SYMBOL_LEN + 1];
-      gfc_symtree *st;
+      gfc_symtree *st = NULL;
 
       m = gfc_match_name (buffer);
       if (m == MATCH_YES)
 	{
-	  st = gfc_find_symtree (gfc_current_ns->sym_root, buffer);
+	  /* Intrinsic functions don't have symtrees yet.  Defer marking
+	     as oacc_functions.  */
+	  if ((isym = gfc_find_function (buffer)) == NULL
+	      && (isym = gfc_find_subroutine (buffer)) == NULL)
+	    st = gfc_find_symtree (gfc_current_ns->sym_root, buffer);
 	  if (st)
 	    {
 	      sym = st->n.sym;
@@ -1730,7 +1745,7 @@ gfc_match_oacc_routine (void)
 	        sym = NULL;
 	    }
 
-	  if (st == NULL
+	  if ((st == NULL && isym == NULL)
 	      || (sym
 		  && !sym->attr.external
 		  && !sym->attr.function
@@ -1764,7 +1779,9 @@ gfc_match_oacc_routine (void)
 	  != MATCH_YES))
     return MATCH_ERROR;
 
-  if (sym != NULL)
+  if (isym != NULL)
+    isym->oacc_function = 1;
+  else if (sym != NULL)
     {
       n = gfc_get_oacc_routine_name ();
       n->sym = sym;
@@ -1782,7 +1799,7 @@ gfc_match_oacc_routine (void)
 				       &old_loc))
 	goto cleanup;
       gfc_current_ns->proc_name->attr.oacc_function
-	= gfc_oacc_routine_dims (c) + 1;
+	= gfc_oacc_routine_dims (c);
     }
 
   if (n)
diff --git a/gcc/fortran/symbol.c b/gcc/fortran/symbol.c
index 0ee7dec..b1dd32b 100644
--- a/gcc/fortran/symbol.c
+++ b/gcc/fortran/symbol.c
@@ -87,6 +87,15 @@ const mstring save_status[] =
     minit ("IMPLICIT-SAVE", SAVE_IMPLICIT),
 };
 
+const mstring oacc_function_types[] =
+{
+  minit ("NONE", OACC_FUNCTION_NONE),
+  minit ("OACC_FUNCTION_SEQ", OACC_FUNCTION_SEQ),
+  minit ("OACC_FUNCTION_GANG", OACC_FUNCTION_GANG),
+  minit ("OACC_FUNCTION_WORKER", OACC_FUNCTION_WORKER),
+  minit ("OACC_FUNCTION_VECTOR", OACC_FUNCTION_VECTOR)
+};
+
 /* This is to make sure the backend generates setup code in the correct
    order.  */
 
diff --git a/gcc/fortran/trans-decl.c b/gcc/fortran/trans-decl.c
index 2f5e434..84fd4ee 100644
--- a/gcc/fortran/trans-decl.c
+++ b/gcc/fortran/trans-decl.c
@@ -1308,30 +1308,34 @@ gfc_add_assign_aux_vars (gfc_symbol * sym)
 }
 
 
-static tree
-add_attributes_to_decl (symbol_attribute sym_attr, tree list)
+tree
+add_omp_offloading_attributes (unsigned omp_declare_target,
+			       enum oacc_function acc_routine, tree list)
 {
-  unsigned id;
-  tree attr;
-
-  for (id = 0; id < EXT_ATTR_NUM; id++)
-    if (sym_attr.ext_attr & (1 << id))
-      {
-	attr = build_tree_list (
-		 get_identifier (ext_attr_list[id].middle_end_name),
-				 NULL_TREE);
-	list = chainon (list, attr);
-      }
-
-  if (sym_attr.omp_declare_target)
+  if (omp_declare_target)
     list = tree_cons (get_identifier ("omp declare target"),
 		      NULL_TREE, list);
 
-  if (sym_attr.oacc_function)
+  if (acc_routine)
     {
       tree dims = NULL_TREE;
       int ix;
-      int level = sym_attr.oacc_function - 1;
+      int level = GOMP_DIM_MAX;
+
+      switch (acc_routine)
+	{
+	case OACC_FUNCTION_GANG:
+	  level = GOMP_DIM_GANG;
+	  break;
+	case OACC_FUNCTION_WORKER:
+	  level = GOMP_DIM_WORKER;
+	  break;
+	case OACC_FUNCTION_VECTOR:
+	  level = GOMP_DIM_VECTOR;
+	  break;
+	case OACC_FUNCTION_SEQ:
+	default:;
+	}
 
       for (ix = GOMP_DIM_MAX; ix--;)
 	dims = tree_cons (build_int_cst (boolean_type_node, ix >= level),
@@ -1344,6 +1348,27 @@ add_attributes_to_decl (symbol_attribute sym_attr, tree list)
   return list;
 }
 
+static tree
+add_attributes_to_decl (symbol_attribute sym_attr, tree list)
+{
+  unsigned id;
+  tree attr;
+
+  for (id = 0; id < EXT_ATTR_NUM; id++)
+    if (sym_attr.ext_attr & (1 << id))
+      {
+	attr = build_tree_list (
+		 get_identifier (ext_attr_list[id].middle_end_name),
+				 NULL_TREE);
+	list = chainon (list, attr);
+      }
+
+  list = add_omp_offloading_attributes (sym_attr.omp_declare_target,
+					sym_attr.oacc_function, list);
+
+  return list;
+}
+
 
 static void build_function_decl (gfc_symbol * sym, bool global);
 
diff --git a/gcc/fortran/trans.c b/gcc/fortran/trans.c
index 28d1341..94eb16d 100644
--- a/gcc/fortran/trans.c
+++ b/gcc/fortran/trans.c
@@ -33,6 +33,7 @@ along with GCC; see the file COPYING3.  If not see
 #include "trans-array.h"
 #include "trans-types.h"
 #include "trans-const.h"
+#include "attribs.h"
 
 /* Naming convention for backend interface code:
 
@@ -2121,11 +2122,18 @@ gfc_unlikely (tree cond, enum br_predictor predictor)
 
   if (optimize)
     {
+      tree fndecl = builtin_decl_explicit (BUILT_IN_EXPECT);
+      tree attributes = NULL_TREE;
+
+      /* Mark calls to BUILT_IN_EXPECT as 'ACC ROUTINE SEQ'.  */
+      if (flag_openacc)
+	attributes = add_omp_offloading_attributes (1, OACC_FUNCTION_SEQ,
+						    attributes);
+
+      decl_attributes (&fndecl, attributes, 0);
       cond = fold_convert (long_integer_type_node, cond);
       tmp = build_zero_cst (long_integer_type_node);
-      cond = build_call_expr_loc (input_location,
-				  builtin_decl_explicit (BUILT_IN_EXPECT),
-				  3, cond, tmp,
+      cond = build_call_expr_loc (input_location, fndecl, 3, cond, tmp,
 				  build_int_cst (integer_type_node,
 						 predictor));
     }
@@ -2143,11 +2151,17 @@ gfc_likely (tree cond, enum br_predictor predictor)
 
   if (optimize)
     {
+      tree fndecl = builtin_decl_explicit (BUILT_IN_EXPECT);
+      tree attributes = NULL_TREE;
+
+      /* Mark calls to BUILT_IN_EXPECT as 'ACC ROUTINE SEQ'.  */
+      if (flag_openacc)
+	attributes = add_omp_offloading_attributes (1, OACC_FUNCTION_SEQ,
+						    attributes);
+      decl_attributes (&fndecl, attributes, 0);
       cond = fold_convert (long_integer_type_node, cond);
       tmp = build_one_cst (long_integer_type_node);
-      cond = build_call_expr_loc (input_location,
-				  builtin_decl_explicit (BUILT_IN_EXPECT),
-				  3, cond, tmp,
+      cond = build_call_expr_loc (input_location, fndecl, 3, cond, tmp,
 				  build_int_cst (integer_type_node,
 						 predictor));
     }
diff --git a/gcc/lto-cgraph.c b/gcc/lto-cgraph.c
index 5cef2ba..552ea6b 100644
--- a/gcc/lto-cgraph.c
+++ b/gcc/lto-cgraph.c
@@ -1201,9 +1201,11 @@ input_overwrite_node (struct lto_file_decl_data *file_data,
 				     LDPR_NUM_KNOWN);
   node->instrumentation_clone = bp_unpack_value (bp, 1);
   node->split_part = bp_unpack_value (bp, 1);
-  gcc_assert (flag_ltrans
-	      || (!node->in_other_partition
-		  && !node->used_from_other_partition));
+
+  int success = flag_ltrans || (!node->in_other_partition
+				&& !node->used_from_other_partition);
+  if (!success)
+    error ("Missing %<%s%>", node->name ());
 }
 
 /* Return string alias is alias of.  */
@@ -1416,9 +1418,11 @@ input_varpool_node (struct lto_file_decl_data *file_data,
     node->set_section_for_node (section);
   node->resolution = streamer_read_enum (ib, ld_plugin_symbol_resolution,
 					        LDPR_NUM_KNOWN);
-  gcc_assert (flag_ltrans
-	      || (!node->in_other_partition
-		  && !node->used_from_other_partition));
+
+  int success = flag_ltrans || (!node->in_other_partition
+				&& !node->used_from_other_partition);
+  if (!success)
+    error ("Missing %<%s%>", node->name ());
 
   return node;
 }
diff --git a/gcc/lto-wrapper.c b/gcc/lto-wrapper.c
index f240812..84b8ad1 100644
--- a/gcc/lto-wrapper.c
+++ b/gcc/lto-wrapper.c
@@ -785,6 +785,8 @@ compile_images_for_offload_targets (unsigned in_argc, char *in_argv[],
 				    struct cl_decoded_option *linker_opts,
 				    unsigned int linker_opt_count)
 {
+  if (!flag_openacc && !flag_openmp && !flag_openmp_simd)
+    return;
   char **names = NULL;
   const char *target_names = getenv (OFFLOAD_TARGET_NAMES_ENV);
   if (!target_names)
@@ -1082,6 +1084,18 @@ run_gcc (unsigned argc, char *argv[])
 	  lto_mode = LTO_MODE_WHOPR;
 	  break;
 
+	case OPT_fopenacc:
+	  flag_openacc = true;
+	  break;
+
+	case OPT_fopenmp:
+	  flag_openmp = true;
+	  break;
+
+	case OPT_fopenmp_simd:
+	  flag_openmp_simd = true;
+	  break;
+
 	default:
 	  break;
 	}
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 22e5909..13e30a6 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -17114,6 +17114,28 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	  default:
 	    break;
 	  }
+      /* Warn if a non-'acc routine' function is called from an OpenACC
+	 offloaded region.  */
+      if (fndecl)
+	{
+	  omp_context *octx = ctx;
+	  bool is_oacc_offloaded = false;
+
+	  /* Check if the current function is an 'acc routine'.  */
+	  if (get_oacc_fn_attrib (current_function_decl) != NULL_TREE)
+	    is_oacc_offloaded = true;
+
+	  while (!is_oacc_offloaded && octx)
+	    {
+	      if (is_oacc_parallel (octx) || is_oacc_kernels (octx))
+		is_oacc_offloaded = true;
+	      octx = octx->outer;
+	    }
+
+	  if (is_oacc_offloaded && get_oacc_fn_attrib (fndecl) == NULL_TREE)
+	    warning_at (gimple_location (call_stmt), 0,
+			"%qE is not an %<acc routine%>", fndecl);
+	}
       /* FALLTHRU */
     default:
       if ((ctx || task_shared_vars)
@@ -19420,7 +19442,8 @@ oacc_loop_fixed_partitions (oacc_loop *loop, unsigned outer_mask)
     {
       unsigned outermost = this_mask & -this_mask;
 
-      if (outermost && outermost <= outer_mask)
+      if ((outermost && outermost <= outer_mask)
+	  || (this_mask && (loop->parent->flags & OLF_SEQ)))
 	{
 	  if (noisy)
 	    {
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-1.c
index 4fcf86e..7afa8c9 100644
--- a/gcc/testsuite/c-c++-common/goacc/kernels-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-1.c
@@ -23,7 +23,7 @@ int
 kernels_noreturn (void)
 {
 #pragma acc kernels
-  __builtin_abort ();
+  __builtin_abort (); /* { dg-warning "'__builtin_abort' is not an 'acc routine'" } */
 
   return 0;
 }
diff --git a/gcc/testsuite/c-c++-common/goacc/parallel-1.c b/gcc/testsuite/c-c++-common/goacc/parallel-1.c
index 6c6cc88..3e070e1 100644
--- a/gcc/testsuite/c-c++-common/goacc/parallel-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/parallel-1.c
@@ -23,7 +23,7 @@ int
 parallel_noreturn (void)
 {
 #pragma acc parallel
-  __builtin_abort ();
+  __builtin_abort (); /* { dg-warning "'__builtin_abort' is not an 'acc routine'" } */
 
   return 0;
 }
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-3.c b/gcc/testsuite/c-c++-common/goacc/routine-3.c
index b322d26..fabae1f 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-3.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-3.c
@@ -49,7 +49,7 @@ main ()
   int red = 0;
 #pragma acc parallel copy (red)
   {
-    /* Independent/seq loop tests.  */
+    /* Independent loop tests.  */
 #pragma acc loop reduction (+:red) // { dg-warning "insufficient partitioning" }
     for (int i = 0; i < 10; i++)
       red += gang ();
@@ -62,6 +62,19 @@ main ()
     for (int i = 0; i < 10; i++)
       red += vector ();
 
+    /* Seq loop tests.  */
+#pragma acc loop seq reduction (+:red)  /* { dg-message "containing loop" } */
+    for (int i = 0; i < 10; i++)
+      red += gang (); /* { dg-error "incorrectly nested" } */
+
+#pragma acc loop seq reduction (+:red)  /* { dg-message "containing loop" } */
+    for (int i = 0; i < 10; i++)
+      red += worker (); /* { dg-error "incorrectly nested" } */
+
+#pragma acc loop seq reduction (+:red)  /* { dg-message "containing loop" } */
+    for (int i = 0; i < 10; i++)
+      red += vector (); /* { dg-error "incorrectly nested" } */
+    
     /* Gang routine tests.  */
 #pragma acc loop gang reduction (+:red)  /* { dg-message "containing loop" } */
     for (int i = 0; i < 10; i++)
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-6.c b/gcc/testsuite/c-c++-common/goacc/routine-6.c
new file mode 100644
index 0000000..fddb5e0
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/routine-6.c
@@ -0,0 +1,26 @@
+/* Test calls to non-routines.  */
+
+int
+sum (int a, int b)
+{
+  return a + b;
+}
+
+int
+main ()
+{
+  int c = 0, i;
+
+#pragma acc parallel loop reduction(+:c)
+  for (i = 0; i < 100; i++)
+    c += sum (i, i); /* { dg-warning "'sum' is not an 'acc routine'" } */
+
+  /* Built-in functions are permitted.  */
+#pragma acc parallel
+  {
+    if (c < 0)
+      __builtin_abort (); /* { dg-warning "'__builtin_abort' is not an 'acc routine'" } */
+  }
+
+  return 0;
+}
diff --git a/gcc/testsuite/gfortran.dg/goacc/fixed-1.f b/gcc/testsuite/gfortran.dg/goacc/fixed-1.f
index 6a454190..0c0fb98 100644
--- a/gcc/testsuite/gfortran.dg/goacc/fixed-1.f
+++ b/gcc/testsuite/gfortran.dg/goacc/fixed-1.f
@@ -1,3 +1,5 @@
+!$ACC ROUTINE(ABORT) SEQ
+
       INTEGER :: ARGC
       ARGC = COMMAND_ARGUMENT_COUNT ()
 
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-7.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-7.f90
new file mode 100644
index 0000000..76b08eb
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-7.f90
@@ -0,0 +1,25 @@
+! Test calls to non-acc routines.
+
+program test
+  implicit none
+  integer c, i
+
+  c = 0
+
+  !$acc parallel loop reduction(+:c)
+  do i = 0, 100
+     c = c + sum (i, i) ! { dg-warning "'sum' is not an 'acc routine'" }
+  end do
+  !$acc end parallel loop
+
+  !$acc parallel
+  if (c .le. 0) call abort ! { dg-warning "is not an 'acc routine'" }
+  !$acc end parallel
+  
+contains
+  integer function sum(a, b)
+    integer a, b
+    sum = a + b
+  end function sum
+  
+end program test
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-8.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-8.f90
new file mode 100644
index 0000000..d2cb51a
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-8.f90
@@ -0,0 +1,122 @@
+! Check routine calls with insufficient parallelism.
+
+! { dg-do compile }
+! { dg-additional-options "-cpp -O0" }
+
+#define M 8
+#define N 32
+
+program main
+  integer :: i
+  integer :: a(N)
+  integer :: b(M * N)
+
+  do i = 1, N
+    a(i) = 0
+  end do
+
+  !$acc parallel copy (a)
+  !$acc loop seq
+    do i = 1, N
+      call seq (a)
+    end do
+  !$acc end parallel
+
+  do i = 1, N
+    if (a(i) .ne.N) call abort
+  end do
+
+  !$acc parallel copy (a)
+  !$acc loop seq ! { dg-message "containing loop here" }
+    do i = 1, N 
+      call gang (a) ! { dg-error "incorrectly nested OpenACC loop parallelism" }
+    end do
+  !$acc end parallel
+
+  do i = 1, N
+    if (a(i) .ne. (N + (N * (-1 * i)))) call abort
+  end do
+
+  do i = 1, N
+    b(i) = i
+  end do
+
+  !$acc parallel copy (b)
+  !$acc loop seq ! { dg-message "containing loop here" }
+    do i = 1, N
+      call worker (b) ! { dg-error "incorrectly nested OpenACC loop parallelism" }
+    end do
+  !$acc end parallel
+
+  do i = 1, N
+    if (b(i) .ne. N + i) call abort
+  end do
+
+  do i = 1, N
+    a(i) = i
+  end do
+
+  !$acc parallel copy (a)
+  !$acc loop seq ! { dg-message "containing loop here" }
+    do i = 1, N
+      call vector (a) ! { dg-error "incorrectly nested OpenACC loop parallelism" }
+    end do
+  !$acc end parallel
+
+  do i = 1, N
+    if (a(i) .ne. 0) call abort
+  end do
+
+contains
+
+subroutine vector (a)
+  !$acc routine vector
+  integer, intent (inout) :: a(N)
+  integer :: i
+
+  !$acc loop vector
+  do i = 1, N
+    a(i) = a(i) - a(i) 
+  end do
+
+end subroutine vector
+
+subroutine worker (b)
+  !$acc routine worker
+  integer, intent (inout) :: b(M*N)
+  integer :: i, j
+
+  !$acc loop worker
+  do i = 1, N
+  !$acc loop vector
+    do j = 1, M
+      b(j + ((i - 1) * M)) = b(j + ((i - 1) * M)) + 1
+    end do
+  end do
+
+end subroutine worker
+
+subroutine gang (a)
+  !$acc routine gang
+  integer, intent (inout) :: a(N)
+  integer :: i
+
+  !$acc loop gang
+  do i = 1, N
+    a(i) = a(i) - i 
+  end do
+
+end subroutine gang
+
+subroutine seq (a)
+  !$acc routine seq
+  integer, intent (inout) :: a(M)
+  integer :: i
+
+  do i = 1, N
+    a(i) = a(i) + 1
+  end do
+
+end subroutine seq
+
+end program main
diff --git a/libgomp/Makefile.am b/libgomp/Makefile.am
index a3e1c2b..c5e7614 100644
--- a/libgomp/Makefile.am
+++ b/libgomp/Makefile.am
@@ -87,8 +87,10 @@ omp_lib_kinds.mod: omp_lib.mod
 	:
 openacc_kinds.mod: openacc.mod
 	:
-openacc.mod: openacc.lo
-	:
+openacc.lo: openacc.f90
+	$(LTFCCOMPILE) -fopenacc -frandom-seed=1 -c -o $@ $^
+openacc.mod: openacc.f90
+	$(FC) $(FCFLAGS) -fopenacc -frandom-seed=1 -c $<
 %.mod: %.f90
 	$(FC) $(FCFLAGS) -fsyntax-only $<
 fortran.lo: libgomp_f.h
diff --git a/libgomp/Makefile.in b/libgomp/Makefile.in
index 88c8517..999409a4 100644
--- a/libgomp/Makefile.in
+++ b/libgomp/Makefile.in
@@ -1286,8 +1286,10 @@ omp_lib_kinds.mod: omp_lib.mod
 	:
 openacc_kinds.mod: openacc.mod
 	:
-openacc.mod: openacc.lo
-	:
+openacc.lo: openacc.f90
+	$(LTFCCOMPILE) -fopenacc -frandom-seed=1 -c -o $@ $^
+openacc.mod: openacc.f90
+	$(FC) $(FCFLAGS) -fopenacc -frandom-seed=1 -c $<
 %.mod: %.f90
 	$(FC) $(FCFLAGS) -fsyntax-only $<
 fortran.lo: libgomp_f.h
diff --git a/libgomp/openacc.f90 b/libgomp/openacc.f90
index 4b71489..98ba493 100644
--- a/libgomp/openacc.f90
+++ b/libgomp/openacc.f90
@@ -128,6 +128,7 @@ module openacc_internal
 
     function acc_on_device_h (d)
       import
+      !$acc routine seq
       integer (acc_device_kind) d
       logical acc_on_device_h
     end function
@@ -719,6 +720,7 @@ end subroutine
 function acc_on_device_h (d)
   use openacc_internal, only: acc_on_device_l
   use openacc_kinds
+  !$acc routine seq
   integer (acc_device_kind) d
   logical acc_on_device_h
   if (acc_on_device_l (d) .eq. 1) then
diff --git a/libgomp/openacc.h b/libgomp/openacc.h
index 7ea8794..e98985c 100644
--- a/libgomp/openacc.h
+++ b/libgomp/openacc.h
@@ -83,6 +83,9 @@ void acc_shutdown (acc_device_t) __GOACC_NOTHROW;
 #ifdef __cplusplus
 int acc_on_device (int __arg) __GOACC_NOTHROW;
 #else
+#ifdef _OPENACC
+#pragma acc routine seq
+#endif
 int acc_on_device (acc_device_t __arg) __GOACC_NOTHROW;
 #endif
 void *acc_malloc (size_t) __GOACC_NOTHROW;
@@ -128,4 +131,8 @@ inline int acc_on_device (acc_device_t __arg) __GOACC_NOTHROW
 }
 #endif
 
+#ifdef _OPENACC
+#pragma acc routine(__builtin_acc_on_device) seq
+#endif
+
 #endif /* _OPENACC_H */
diff --git a/libgomp/openacc_lib.h b/libgomp/openacc_lib.h
index a3f94d7..d627857 100644
--- a/libgomp/openacc_lib.h
+++ b/libgomp/openacc_lib.h
@@ -142,6 +142,7 @@
       interface acc_on_device
         function acc_on_device_h (devicetype)
           import acc_device_kind
+!$acc routine seq
           logical acc_on_device_h
           integer (acc_device_kind) devicetype
         end function
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/abort-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/abort-1.c
index 296708f..bc4eab3 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/abort-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/abort-1.c
@@ -2,6 +2,7 @@
 
 #include <stdio.h>
 #include <stdlib.h>
+#pragma acc routine(abort) seq
 
 int
 main (void)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/abort-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/abort-2.c
index debb81e..20076cd 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/abort-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/abort-2.c
@@ -1,6 +1,7 @@
 /* { dg-do run } */
 
 #include <stdlib.h>
+#pragma acc routine(abort) seq
 
 int
 main (int argc, char **argv)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/abort-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/abort-3.c
index bca425e..e6fc72f 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/abort-3.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/abort-3.c
@@ -2,6 +2,7 @@
 
 #include <stdio.h>
 #include <stdlib.h>
+#pragma acc routine(abort) seq
 
 int
 main (void)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/abort-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/abort-4.c
index c29ca3f..53a069a 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/abort-4.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/abort-4.c
@@ -1,6 +1,7 @@
 /* { dg-do run } */
 
 #include <stdlib.h>
+#pragma acc routine(abort) seq
 
 int
 main (int argc, char **argv)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/abort-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/abort-5.c
index 314f04a..c38576e 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/abort-5.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/abort-5.c
@@ -2,6 +2,7 @@
 /* { dg-additional-options "-flto" { target lto } } */
 
 #include <stdlib.h>
+#pragma acc routine(abort) seq
 
 int
 main (int argc, char **argv)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_on_device-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_on_device-1.c
index 8112745..a214329 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_on_device-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_on_device-1.c
@@ -4,6 +4,7 @@
 
 #include <stdlib.h>
 #include <openacc.h>
+#pragma acc routine(abort) seq
 
 int
 main (int argc, char *argv[])
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-clauses-kernels-ipa-pta.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-clauses-kernels-ipa-pta.c
index 2cd98bd..83d8e56 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-clauses-kernels-ipa-pta.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-clauses-kernels-ipa-pta.c
@@ -1,4 +1,4 @@
 /* { dg-do run { target lto } } */
-/* { dg-additional-options "-fipa-pta -flto -flto-partition=max" } */
+/* { dg-additional-options "-fipa-pta -flto -flto-partition=max -fno-exceptions" } */
 
 #include "data-clauses-kernels.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-clauses-kernels.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-clauses-kernels.c
index f7f2d1c..a3934cb 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-clauses-kernels.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-clauses-kernels.c
@@ -1,2 +1,4 @@
+/* { dg-additional-options "-fno-exceptions" }  */
+
 #define CONSTRUCT kernels
 #include "data-clauses.h"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-clauses-parallel-ipa-pta.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-clauses-parallel-ipa-pta.c
index ddcf4e3..6d24b3d 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-clauses-parallel-ipa-pta.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-clauses-parallel-ipa-pta.c
@@ -1,4 +1,4 @@
 /* { dg-do run { target lto } } */
-/* { dg-additional-options "-fipa-pta -flto -flto-partition=max" } */
+/* { dg-additional-options "-fipa-pta -flto -flto-partition=max -fno-exceptions" } */
 
 #include "data-clauses-parallel.c"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-clauses-parallel.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-clauses-parallel.c
index e734b2f..02f1e88 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-clauses-parallel.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-clauses-parallel.c
@@ -1,2 +1,4 @@
+/* { dg-additional-options "-fno-exceptions" }  */
+
 #define CONSTRUCT parallel
 #include "data-clauses.h"
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-clauses.h b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-clauses.h
index d557bef..5e7eb14 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-clauses.h
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-clauses.h
@@ -1,3 +1,5 @@
+#pragma acc routine(__builtin_abort) seq
+
 int i;
 
 int main(void)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c
index 5398905..81aec7e 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c
@@ -2,6 +2,8 @@
 #include <stdlib.h>
 #include <stdbool.h>
 
+#pragma acc routine(abort) seq
+
 #define N   32
 
 int
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-1.c
index 7bff6cd..21f8bc1 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-1.c
@@ -3,6 +3,7 @@
    { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */
 
 #include <stdio.h>
+#include <openacc.h>
 
 #define N (32*32*32+17)
 int main ()
@@ -20,7 +21,7 @@ int main ()
 #pragma acc loop gang
     for (unsigned ix = 0; ix < N; ix++)
       {
-	if (__builtin_acc_on_device (5))
+	if (__builtin_acc_on_device (acc_device_nvidia))
 	  {
 	    int g = 0, w = 0, v = 0;
 
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-2.c
index 92b82a0..72c3bde 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-2.c
@@ -3,6 +3,7 @@
    { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */
 
 #include <stdio.h>
+#include <openacc.h>
 
 #define N (32*32*32+17)
 int main ()
@@ -20,7 +21,7 @@ int main ()
 #pragma acc loop gang (static:1)
     for (unsigned ix = 0; ix < N; ix++)
       {
-	if (__builtin_acc_on_device (5))
+	if (__builtin_acc_on_device (acc_device_nvidia))
 	  {
 	    int g = 0, w = 0, v = 0;
 
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c
index 42b612a..364f058 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c
@@ -3,6 +3,7 @@
    { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */
 
 #include <stdio.h>
+#include <openacc.h>
 
 #define N (32*32*32+17)
 int main ()
@@ -20,7 +21,7 @@ int main ()
 #pragma acc loop gang worker vector
     for (unsigned ix = 0; ix < N; ix++)
       {
-	if (__builtin_acc_on_device (5))
+	if (__builtin_acc_on_device (acc_device_nvidia))
 	  {
 	    int g = 0, w = 0, v = 0;
 
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c
index a8684f95..d1d27b3 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c
@@ -3,6 +3,7 @@
    { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */
 
 #include <stdio.h>
+#include <openacc.h>
 
 #define N (32*32*32+17)
 int main ()
@@ -18,7 +19,7 @@ int main ()
       {
 	int val = ix;
 	
-	if (__builtin_acc_on_device (5))
+	if (__builtin_acc_on_device (acc_device_nvidia))
 	  {
 	    int g = 0, w = 0, v = 0;
 
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-gwv-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-gwv-1.c
index 3b104cf..0ebbc63 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-gwv-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-gwv-1.c
@@ -3,6 +3,7 @@
    { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */
 
 #include <stdio.h>
+#include <openacc.h>
 
 #define N (32*32*32+17)
 int main ()
@@ -18,7 +19,7 @@ int main ()
       {
 	int val = ix;
 	
-	if (__builtin_acc_on_device (5))
+	if (__builtin_acc_on_device (acc_device_nvidia))
 	  {
 	    int g = 0, w = 0, v = 0;
 
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-1.c
index b77ae76..1b350e9 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-1.c
@@ -3,6 +3,7 @@
    { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */
 
 #include <stdio.h>
+#include <openacc.h>
 
 #define N (32*32*32+17)
 
@@ -19,7 +20,7 @@ int main ()
       {
 	int val = ix;
 	
-	if (__builtin_acc_on_device (5))
+	if (__builtin_acc_on_device (acc_device_nvidia))
 	  {
 	    int g = 0, w = 0, v = 0;
 
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c
index 16d8f9f..4b6d835 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c
@@ -3,6 +3,7 @@
    { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */
 
 #include <stdio.h>
+#include <openacc.h>
 
 #define N (32*32*32+17)
 
@@ -21,7 +22,7 @@ int main ()
       {
 	int val = ix;
 	
-	if (__builtin_acc_on_device (5))
+	if (__builtin_acc_on_device (acc_device_nvidia))
 	  {
 	    int g = 0, w = 0, v = 0;
 
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c
index 9cc12b3..44ab546 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c
@@ -3,6 +3,7 @@
    { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */
 
 #include <stdio.h>
+#include <openacc.h>
 
 #define N (32*32*32+17)
 int main ()
@@ -18,7 +19,7 @@ int main ()
       {
 	int val = ix;
 	
-	if (__builtin_acc_on_device (5))
+	if (__builtin_acc_on_device (acc_device_nvidia))
 	  {
 	    int g = 0, w = 0, v = 0;
 
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c
index f0c9d81..2e3f1e5 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c
@@ -3,6 +3,7 @@
    { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */
 
 #include <stdio.h>
+#include <openacc.h>
 
 #define N (32*32*32+17)
 int main ()
@@ -20,7 +21,7 @@ int main ()
       {
 	int val = ix;
 	
-	if (__builtin_acc_on_device (5))
+	if (__builtin_acc_on_device (acc_device_nvidia))
 	  {
 	    int g = 0, w = 0, v = 0;
 
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-wv-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-wv-1.c
index 398b7cc..30f767a 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-wv-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-wv-1.c
@@ -2,6 +2,7 @@
 /* { dg-additional-options "-O2" } */
 
 #include <stdio.h>
+#include <openacc.h>
 
 #define N (32*32*32+17)
 int main ()
@@ -17,7 +18,7 @@ int main ()
       {
 	int val = ix;
 	
-	if (__builtin_acc_on_device (5))
+	if (__builtin_acc_on_device (acc_device_nvidia))
 	  {
 	    int g = 0, w = 0, v = 0;
 
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-v-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-v-1.c
index 2974807..7a0d688 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-v-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-v-1.c
@@ -3,6 +3,7 @@
    { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */
 
 #include <stdio.h>
+#include <openacc.h>
 
 #define N (32*32*32+17)
 int main ()
@@ -20,7 +21,7 @@ int main ()
 #pragma acc loop vector
     for (unsigned ix = 0; ix < N; ix++)
       {
-	if (__builtin_acc_on_device (5))
+	if (__builtin_acc_on_device (acc_device_nvidia))
 	  {
 	    int g = 0, w = 0, v = 0;
 
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c
index 33b6eae..c165f1d 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c
@@ -3,6 +3,7 @@
    { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */
 
 #include <stdio.h>
+#include <openacc.h>
 
 #define N (32*32*32+17)
 int main ()
@@ -20,7 +21,7 @@ int main ()
 #pragma acc loop worker
     for (unsigned ix = 0; ix < N; ix++)
       {
-	if (__builtin_acc_on_device (5))
+	if (__builtin_acc_on_device (acc_device_nvidia))
 	  {
 	    int g = 0, w = 0, v = 0;
 
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c
index 578cfad..70bfa62 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c
@@ -3,6 +3,7 @@
    { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */
 
 #include <stdio.h>
+#include <openacc.h>
 
 #define N (32*32*32+17)
 int main ()
@@ -20,7 +21,7 @@ int main ()
 #pragma acc loop worker vector
     for (unsigned ix = 0; ix < N; ix++)
       {
-	if (__builtin_acc_on_device (5))
+	if (__builtin_acc_on_device (acc_device_nvidia))
 	  {
 	    int g = 0, w = 0, v = 0;
 
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-4.c
index d6ff44d..02b1f15 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-4.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-4.c
@@ -78,7 +78,7 @@ main(int argc, char **argv)
 
 #pragma acc parallel copy (a[0:N])
   {
-#pragma acc loop seq
+#pragma acc loop /* { dg-warning "insufficient partitioning" } */
     for (i = 0; i < N; i++)
       gang (&a[0]);
   }
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c
index 9d14c3b..be80457 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c
@@ -3,6 +3,7 @@
    { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */
 
 #include <stdio.h>
+#include <openacc.h>
 
 #define N (32*32*32+17)
 
@@ -12,7 +13,7 @@ void __attribute__ ((noinline)) gang (int ary[N])
 #pragma acc loop gang
     for (unsigned ix = 0; ix < N; ix++)
       {
-	if (__builtin_acc_on_device (5))
+	if (__builtin_acc_on_device (acc_device_nvidia))
 	  {
 	    int g = 0, w = 0, v = 0;
 
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c
index ace2f49..b6c689b 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c
@@ -3,6 +3,7 @@
    { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */
 
 #include <stdio.h>
+#include <openacc.h>
 
 #define N (32*32*32+17)
 
@@ -12,7 +13,7 @@ void __attribute__ ((noinline)) gang (int ary[N])
 #pragma acc loop gang worker vector
   for (unsigned ix = 0; ix < N; ix++)
     {
-      if (__builtin_acc_on_device (5))
+      if (__builtin_acc_on_device (acc_device_nvidia))
 	{
 	  int g = 0, w = 0, v = 0;
 
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-v-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-v-1.c
index 2503e8d..5a73b0b 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-v-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-v-1.c
@@ -3,6 +3,7 @@
    { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */
 
 #include <stdio.h>
+#include <openacc.h>
 
 #define N (32*32*32+17)
 
@@ -12,7 +13,7 @@ void __attribute__ ((noinline)) vector (int ary[N])
 #pragma acc loop vector
   for (unsigned ix = 0; ix < N; ix++)
     {
-      if (__builtin_acc_on_device (5))
+      if (__builtin_acc_on_device (acc_device_nvidia))
 	{
 	  int g = 0, w = 0, v = 0;
 
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c
index 80cd462..523353a 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c
@@ -3,6 +3,7 @@
    { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */
 
 #include <stdio.h>
+#include <openacc.h>
 
 #define N (32*32*32+17)
 
@@ -12,7 +13,7 @@ void __attribute__ ((noinline)) worker (int ary[N])
 #pragma acc loop worker
   for (unsigned ix = 0; ix < N; ix++)
     {
-      if (__builtin_acc_on_device (5))
+      if (__builtin_acc_on_device (acc_device_nvidia))
 	{
 	  int g = 0, w = 0, v = 0;
 
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c
index 5e45fad..e92b160 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c
@@ -3,6 +3,7 @@
    { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */
 
 #include <stdio.h>
+#include <openacc.h>
 
 #define N (32*32*32+17)
 
@@ -12,7 +13,7 @@ void __attribute__ ((noinline)) worker (int ary[N])
 #pragma acc loop worker vector
   for (unsigned ix = 0; ix < N; ix++)
     {
-      if (__builtin_acc_on_device (5))
+      if (__builtin_acc_on_device (acc_device_nvidia))
 	{
 	  int g = 0, w = 0, v = 0;
 
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-type-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-type-1.c
index 5adfcec..6c2b1c2 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-type-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-type-1.c
@@ -1,3 +1,5 @@
+#pragma acc routine(__builtin_abort) seq
+
 #define vector __attribute__ ((vector_size (4 * sizeof(int))))
 
 int main(void)
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/abort-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/abort-1.f90
index b38303d..48ebc38 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/abort-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/abort-1.f90
@@ -1,5 +1,6 @@
 program main
   implicit none
+  !$acc routine(abort) seq
 
   print *, "CheCKpOInT"
   !$acc parallel
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/abort-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/abort-2.f90
index 2ba2bcb..a80593e 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/abort-2.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/abort-2.f90
@@ -1,5 +1,6 @@
 program main
   implicit none
+  !$acc routine(abort) seq
 
   integer :: argc
   argc = command_argument_count ()
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90
index 1a10f32..94e45b3 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90
@@ -6,6 +6,7 @@
 
 use openacc
 implicit none
+!$acc routine(abort) seq
 
 ! Host.
 
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f b/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f
index a19045b..cbd1dd9 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f
@@ -6,6 +6,7 @@
 
       USE OPENACC
       IMPLICIT NONE
+!$ACC ROUTINE(ABORT) SEQ
 
 !Host.
 
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f b/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f
index c391776..3e016f4 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f
@@ -6,6 +6,7 @@
 
       IMPLICIT NONE
       INCLUDE "openacc_lib.h"
+!$ACC ROUTINE(ABORT) SEQ
 
 !Host.
 
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/nested-function-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/nested-function-1.f90
index fdbca44..2b14159 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/nested-function-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/nested-function-1.f90
@@ -3,6 +3,8 @@
 ! { dg-do run }
 
 program collapse2
+  !$acc routine(abort) seq
+
   call test1
   call test2
 contains
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90
index 200188e..27cda44 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90
@@ -1,121 +1,101 @@
+! Test acc routines.
 
 ! { dg-do run }
-! { dg-additional-options "-cpp" }
 
-#define M 8
-#define N 32
+module size
+  integer, parameter :: N = 32
+end module size
 
 program main
+  use size
+  implicit none
+
   integer :: i
   integer :: a(N)
-  integer :: b(M * N)
 
-  do i = 1, N
-    a(i) = 0
-  end do
-
-  !$acc parallel copy (a)
-  !$acc loop seq
-    do i = 1, N
-      call seq (a)
-    end do
+  !$acc parallel
+  call seq (a)
   !$acc end parallel
 
   do i = 1, N
-    if (a(i) .ne.N) call abort
+    if (a(i) .ne. 4) call abort
   end do
 
-  !$acc parallel copy (a)
-  !$acc loop seq
-    do i = 1, N 
-      call gang (a)
-    end do
+  !$acc parallel
+  call gang (a)
   !$acc end parallel
 
   do i = 1, N
-    if (a(i) .ne. (N + (N * (-1 * i)))) call abort
+    if (a(i) .ne. 3) call abort
   end do
 
-  do i = 1, N
-    b(i) = i
-  end do
-
-  !$acc parallel copy (b)
-  !$acc loop seq
-    do i = 1, N
-      call worker (b)
-    end do
+  !$acc parallel
+  call worker (a)
   !$acc end parallel
 
   do i = 1, N
-    if (b(i) .ne. N + i) call abort
-  end do
-
-  do i = 1, N
-    a(i) = i
+    if (a(i) .ne. 2) call abort
   end do
 
-  !$acc parallel copy (a)
-  !$acc loop seq
-    do i = 1, N
-      call vector (a)
-    end do
+  !$acc parallel
+  call vector (a)
   !$acc end parallel
 
   do i = 1, N
-    if (a(i) .ne. 0) call abort
+    if (a(i) .ne. 1) call abort
   end do
 
 contains
 
 subroutine vector (a)
+  use size
+  implicit none
   !$acc routine vector
   integer, intent (inout) :: a(N)
   integer :: i
 
   !$acc loop vector
   do i = 1, N
-    a(i) = a(i) - a(i) 
+    a(i) = 1
   end do
-
 end subroutine vector
 
-subroutine worker (b)
+subroutine worker (a)
+  use size
+  implicit none
   !$acc routine worker
-  integer, intent (inout) :: b(M*N)
-  integer :: i, j
+  integer, intent (inout) :: a(N)
+  integer :: i
 
   !$acc loop worker
   do i = 1, N
-  !$acc loop vector
-    do j = 1, M
-      b(j + ((i - 1) * M)) = b(j + ((i - 1) * M)) + 1
-    end do
+    a(i) = 2
   end do
-
 end subroutine worker
 
 subroutine gang (a)
+  use size
+  implicit none
   !$acc routine gang
   integer, intent (inout) :: a(N)
   integer :: i
 
   !$acc loop gang
   do i = 1, N
-    a(i) = a(i) - i 
+    a(i) = 3
   end do
-
 end subroutine gang
 
 subroutine seq (a)
+  use size
+  implicit none
   !$acc routine seq
-  integer, intent (inout) :: a(M)
+  integer, intent (inout) :: a(N)
   integer :: i
 
   do i = 1, N
-    a(i) = a(i) + 1
+    a(i) = 4
   end do
-
 end subroutine seq
 
 end program main

Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]