diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index e3899bae505..8c8722ccfbe 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -15437,6 +15437,41 @@ c_parser_oacc_clause_wait (c_parser *parser, tree list)
   return list;
 }
 
+/* OpenACC 2.7:
+   self [( expression )] */
+
+static tree
+c_parser_oacc_compute_clause_self (c_parser *parser, tree list)
+{
+  tree t;
+  location_t location = c_parser_peek_token (parser)->location;
+  if (c_parser_peek_token (parser)->type == CPP_OPEN_PAREN)
+    {
+      matching_parens parens;
+      parens.consume_open (parser);
+
+      location_t loc = c_parser_peek_token (parser)->location;
+      c_expr expr = c_parser_expr_no_commas (parser, NULL);
+      expr = convert_lvalue_to_rvalue (loc, expr, true, true);
+      t = c_objc_common_truthvalue_conversion (loc, expr.value);
+      t = c_fully_fold (t, false, NULL);
+      parens.skip_until_found_close (parser);
+    }
+  else
+    t = truthvalue_true_node;
+
+  for (tree c = list; c; c = OMP_CLAUSE_CHAIN (c))
+    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_SELF)
+      {
+	error_at (location, "too many %<self%> clauses");
+	return list;
+      }
+
+  tree c = build_omp_clause (location, OMP_CLAUSE_SELF);
+  OMP_CLAUSE_SELF_EXPR (c) = t;
+  OMP_CLAUSE_CHAIN (c) = list;
+  return c;
+}
 
 /* OpenMP 5.0:
    order ( concurrent )
@@ -17560,7 +17595,8 @@ c_parser_omp_clause_detach (c_parser *parser, tree list)
 
 static tree
 c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
-			   const char *where, bool finish_p = true)
+			   const char *where, bool finish_p = true,
+			   bool compute_p = false)
 {
   tree clauses = NULL;
   bool first = true;
@@ -17576,7 +17612,18 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
 	c_parser_consume_token (parser);
 
       here = c_parser_peek_token (parser)->location;
-      c_kind = c_parser_omp_clause_name (parser);
+
+      /* For OpenACC compute directives */
+      if (compute_p
+	  && c_parser_next_token_is (parser, CPP_NAME)
+	  && !strcmp (IDENTIFIER_POINTER (c_parser_peek_token (parser)->value),
+		      "self"))
+	{
+	  c_kind = PRAGMA_OACC_CLAUSE_SELF;
+	  c_parser_consume_token (parser);
+	}
+      else
+	c_kind = c_parser_omp_clause_name (parser);
 
       switch (c_kind)
 	{
@@ -17708,6 +17755,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
 					     false, clauses);
 	  c_name = "reduction";
 	  break;
+	case PRAGMA_OACC_CLAUSE_SELF:
+	  clauses = c_parser_oacc_compute_clause_self (parser, clauses);
+	  c_name = "self";
+	  break;
 	case PRAGMA_OACC_CLAUSE_SEQ:
 	  clauses = c_parser_oacc_simple_clause (here, OMP_CLAUSE_SEQ,
 						 clauses);
@@ -18544,6 +18595,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SELF)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
 
@@ -18564,6 +18616,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SELF)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
 
@@ -18582,6 +18635,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SELF)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
 
 static tree
@@ -18624,7 +18678,7 @@ c_parser_oacc_compute (location_t loc, c_parser *parser,
 	}
     }
 
-  tree clauses = c_parser_oacc_all_clauses (parser, mask, p_name);
+  tree clauses = c_parser_oacc_all_clauses (parser, mask, p_name, true, true);
 
   tree block = c_begin_omp_parallel ();
   add_stmt (c_parser_omp_structured_block (parser, if_p));
diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc
index 22e240a3c2a..89bf22df372 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -15821,6 +15821,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	  continue;
 
 	case OMP_CLAUSE_IF:
+	case OMP_CLAUSE_SELF:
 	case OMP_CLAUSE_NUM_THREADS:
 	case OMP_CLAUSE_NUM_TEAMS:
 	case OMP_CLAUSE_THREAD_LIMIT:
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index 20dbe1089ec..efa805fb3b0 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -40956,13 +40956,51 @@ cp_parser_oacc_clause_async (cp_parser *parser, tree list)
   return list;
 }
 
+/* OpenACC 2.7:
+   self [( expression )] */
+
+static tree
+cp_parser_oacc_compute_clause_self (cp_parser *parser, tree list)
+{
+  tree t;
+  location_t location = cp_lexer_peek_token (parser->lexer)->location;
+  if (cp_lexer_peek_token (parser->lexer)->type == CPP_OPEN_PAREN)
+    {
+      matching_parens parens;
+      parens.consume_open (parser);
+      t = cp_parser_assignment_expression (parser);
+      if (t == error_mark_node
+	  || !parens.require_close (parser))
+	{
+	  cp_parser_skip_to_closing_parenthesis (parser, /*recovering=*/true,
+						 /*or_comma=*/false,
+						 /*consume_paren=*/true);
+	  return list;
+	}
+    }
+  else
+    t = truthvalue_true_node;
+
+  for (tree c = list; c; c = OMP_CLAUSE_CHAIN (c))
+    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_SELF)
+      {
+	error_at (location, "too many %<self%> clauses");
+	return list;
+      }
+
+  tree c = build_omp_clause (location, OMP_CLAUSE_SELF);
+  OMP_CLAUSE_SELF_EXPR (c) = t;
+  OMP_CLAUSE_CHAIN (c) = list;
+  return c;
+}
+
 /* Parse all OpenACC clauses.  The set clauses allowed by the directive
    is a bitmask in MASK.  Return the list of clauses found.  */
 
 static tree
 cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
 			    const char *where, cp_token *pragma_tok,
-			    bool finish_p = true)
+			    bool finish_p = true, bool compute_p = false)
 {
   tree clauses = NULL;
   bool first = true;
@@ -40982,7 +41020,19 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
 	cp_lexer_consume_token (parser->lexer);
 
       here = cp_lexer_peek_token (parser->lexer)->location;
-      c_kind = cp_parser_omp_clause_name (parser);
+
+      /* For OpenACC compute directives */
+      if (compute_p
+	  && cp_lexer_next_token_is (parser->lexer, CPP_NAME)
+	  && !strcmp (IDENTIFIER_POINTER
+		      (cp_lexer_peek_token (parser->lexer)->u.value),
+		      "self"))
+	{
+	  c_kind = PRAGMA_OACC_CLAUSE_SELF;
+	  cp_lexer_consume_token (parser->lexer);
+	}
+      else
+	c_kind = cp_parser_omp_clause_name (parser);
 
       switch (c_kind)
 	{
@@ -41116,6 +41166,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
 					      false, clauses);
 	  c_name = "reduction";
 	  break;
+	case PRAGMA_OACC_CLAUSE_SELF:
+	  clauses = cp_parser_oacc_compute_clause_self (parser, clauses);
+	  c_name = "self";
+	  break;
 	case PRAGMA_OACC_CLAUSE_SEQ:
 	  clauses = cp_parser_oacc_simple_clause (here, OMP_CLAUSE_SEQ,
 						  clauses);
@@ -46227,6 +46281,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SELF)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
 
@@ -46247,6 +46302,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SELF)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH)       \
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
 
@@ -46265,6 +46321,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SELF)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
 
 static tree
@@ -46310,7 +46367,8 @@ cp_parser_oacc_compute (cp_parser *parser, cp_token *pragma_tok,
 	}
     }
 
-  tree clauses = cp_parser_oacc_all_clauses (parser, mask, p_name, pragma_tok);
+  tree clauses = cp_parser_oacc_all_clauses (parser, mask, p_name, pragma_tok,
+					     true, true);
 
   tree block = begin_omp_parallel ();
   unsigned int save = cp_parser_begin_omp_structured_block (parser);
diff --git a/gcc/cp/pt.cc b/gcc/cp/pt.cc
index 6b20c58ce66..f8d8fc31c53 100644
--- a/gcc/cp/pt.cc
+++ b/gcc/cp/pt.cc
@@ -18037,6 +18037,7 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort,
 	  /* FALLTHRU */
 	case OMP_CLAUSE_TILE:
 	case OMP_CLAUSE_IF:
+	case OMP_CLAUSE_SELF:
 	case OMP_CLAUSE_NUM_THREADS:
 	case OMP_CLAUSE_SCHEDULE:
 	case OMP_CLAUSE_COLLAPSE:
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index a2e74a5d2c7..ddb9f10b45a 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -7344,13 +7344,14 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	  goto handle_field_decl;
 
 	case OMP_CLAUSE_IF:
-	  t = OMP_CLAUSE_IF_EXPR (c);
+	case OMP_CLAUSE_SELF:
+	  t = OMP_CLAUSE_OPERAND (c, 0);
 	  t = maybe_convert_cond (t);
 	  if (t == error_mark_node)
 	    remove = true;
 	  else if (!processing_template_decl)
 	    t = fold_build_cleanup_point_expr (TREE_TYPE (t), t);
-	  OMP_CLAUSE_IF_EXPR (c) = t;
+	  OMP_CLAUSE_OPERAND (c, 0) = t;
 	  break;
 
 	case OMP_CLAUSE_FINAL:
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index a58c60e9828..7a21a6cb008 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1540,6 +1540,7 @@ typedef struct gfc_omp_clauses
 {
   gfc_omp_namelist *lists[OMP_LIST_NUM];
   struct gfc_expr *if_expr;
+  struct gfc_expr *self_expr;
   struct gfc_expr *final_expr;
   struct gfc_expr *num_threads;
   struct gfc_expr *chunk_size;
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index 9714d5f8f61..038907baa48 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -1091,6 +1091,7 @@ enum omp_mask2
   OMP_CLAUSE_ENTER, /* OpenMP 5.2 */
   OMP_CLAUSE_DOACROSS, /* OpenMP 5.2 */
   OMP_CLAUSE_ASSUMPTIONS, /* OpenMP 5.1. */
+  OMP_CLAUSE_SELF, /* OpenACC 2.7  */
   /* This must come last.  */
   OMP_MASK2_LAST
 };
@@ -3412,6 +3413,27 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 	      else
 		gfc_current_locus = old_loc;
 	    }
+	  if ((mask & OMP_CLAUSE_SELF)
+	      && (m = gfc_match_dupl_check (!c->self_expr, "self"))
+	          != MATCH_NO)
+	    {
+	      gcc_assert (!(mask & OMP_CLAUSE_HOST_SELF));
+	      if (m == MATCH_ERROR)
+		goto error;
+	      m = gfc_match (" ( %e )", &c->self_expr);
+	      if (m == MATCH_ERROR)
+		{
+		  gfc_current_locus = old_loc;
+		  break;
+		}
+	      else if (m == MATCH_NO)
+		{
+		  c->self_expr = gfc_get_logical_expr (gfc_default_logical_kind,
+						       NULL, true);
+		  needs_space = true;
+		}
+	      continue;
+	    }
 	  if ((mask & OMP_CLAUSE_HOST_SELF)
 	      && gfc_match ("self ( ") == MATCH_YES
 	      && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
@@ -3677,19 +3699,22 @@ error:
    | OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT		      \
    | OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT	      \
    | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE      \
-   | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT | OMP_CLAUSE_ATTACH)
+   | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT | OMP_CLAUSE_ATTACH		      \
+   | OMP_CLAUSE_SELF)
 #define OACC_KERNELS_CLAUSES \
   (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_NUM_GANGS	      \
    | OMP_CLAUSE_NUM_WORKERS | OMP_CLAUSE_VECTOR_LENGTH | OMP_CLAUSE_DEVICEPTR \
    | OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT		      \
    | OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT	      \
-   | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT | OMP_CLAUSE_ATTACH)
+   | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT | OMP_CLAUSE_ATTACH	              \
+   | OMP_CLAUSE_SELF)
 #define OACC_SERIAL_CLAUSES \
   (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_REDUCTION	      \
    | OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT		      \
    | OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT	      \
    | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE      \
-   | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT | OMP_CLAUSE_ATTACH)
+   | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT | OMP_CLAUSE_ATTACH	              \
+   | OMP_CLAUSE_SELF)
 #define OACC_DATA_CLAUSES \
   (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_DEVICEPTR  | OMP_CLAUSE_COPY	      \
    | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT | OMP_CLAUSE_CREATE		      \
@@ -7251,6 +7276,14 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 		   &expr->where);
       if_without_mod = true;
     }
+  if (omp_clauses->self_expr)
+    {
+      gfc_expr *expr = omp_clauses->self_expr;
+      if (!gfc_resolve_expr (expr)
+	  || expr->ts.type != BT_LOGICAL || expr->rank != 0)
+	gfc_error ("SELF clause at %L requires a scalar LOGICAL expression",
+		   &expr->where);
+    }
   for (ifc = 0; ifc < OMP_IF_LAST; ifc++)
     if (omp_clauses->if_exprs[ifc])
       {
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index 340da85cb99..0f8323901d7 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -3943,6 +3943,22 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
       OMP_CLAUSE_IF_EXPR (c) = if_var;
       omp_clauses = gfc_trans_add_clause (c, omp_clauses);
     }
+
+  if (clauses->self_expr)
+    {
+      tree self_var;
+
+      gfc_init_se (&se, NULL);
+      gfc_conv_expr (&se, clauses->self_expr);
+      gfc_add_block_to_block (block, &se.pre);
+      self_var = gfc_evaluate_now (se.expr, block);
+      gfc_add_block_to_block (block, &se.post);
+
+      c = build_omp_clause (gfc_get_location (&where), OMP_CLAUSE_SELF);
+      OMP_CLAUSE_SELF_EXPR (c) = self_var;
+      omp_clauses = gfc_trans_add_clause (c, omp_clauses);
+    }
+
   for (ifc = 0; ifc < OMP_IF_LAST; ifc++)
     if (clauses->if_exprs[ifc])
       {
@@ -6595,6 +6611,8 @@ gfc_split_omp_clauses (gfc_code *code,
 	  /* And this is copied to all.  */
 	  clausesa[GFC_OMP_SPLIT_TARGET].if_expr
 	    = code->ext.omp_clauses->if_expr;
+	  clausesa[GFC_OMP_SPLIT_TARGET].self_expr
+	    = code->ext.omp_clauses->self_expr;
 	  clausesa[GFC_OMP_SPLIT_TARGET].nowait
 	    = code->ext.omp_clauses->nowait;
 	}
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 6f1e95bfbe7..d309459be1b 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -11867,6 +11867,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	    }
 	  /* Fall through.  */
 
+	case OMP_CLAUSE_SELF:
 	case OMP_CLAUSE_FINAL:
 	  OMP_CLAUSE_OPERAND (c, 0)
 	    = gimple_boolify (OMP_CLAUSE_OPERAND (c, 0));
@@ -13093,6 +13094,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 	case OMP_CLAUSE_COPYIN:
 	case OMP_CLAUSE_COPYPRIVATE:
 	case OMP_CLAUSE_IF:
+	case OMP_CLAUSE_SELF:
 	case OMP_CLAUSE_NUM_THREADS:
 	case OMP_CLAUSE_NUM_TEAMS:
 	case OMP_CLAUSE_THREAD_LIMIT:
diff --git a/gcc/omp-expand.cc b/gcc/omp-expand.cc
index db58b3cb49b..f7b7c67d459 100644
--- a/gcc/omp-expand.cc
+++ b/gcc/omp-expand.cc
@@ -10331,6 +10331,47 @@ expand_omp_target (struct omp_region *region)
 	}
     }
 
+  if ((c = omp_find_clause (clauses, OMP_CLAUSE_SELF)) != NULL_TREE)
+    {
+      gcc_assert (is_gimple_omp_oacc (entry_stmt) && offloaded);
+
+      edge e = split_block_after_labels (new_bb);
+      basic_block cond_bb = e->src;
+      new_bb = e->dest;
+      remove_edge (e);
+
+      basic_block then_bb = create_empty_bb (cond_bb);
+      basic_block else_bb = create_empty_bb (then_bb);
+      set_immediate_dominator (CDI_DOMINATORS, then_bb, cond_bb);
+      set_immediate_dominator (CDI_DOMINATORS, else_bb, cond_bb);
+
+      tree self_cond = gimple_boolify (OMP_CLAUSE_SELF_EXPR (c));
+      stmt = gimple_build_cond_empty (self_cond);
+      gsi = gsi_last_bb (cond_bb);
+      gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
+
+      tree tmp_var = create_tmp_var (TREE_TYPE (goacc_flags));
+      stmt = gimple_build_assign (tmp_var, BIT_IOR_EXPR, goacc_flags,
+				  build_int_cst (integer_type_node,
+						 GOACC_FLAG_LOCAL_DEVICE));
+      gsi = gsi_start_bb (then_bb);
+      gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
+
+      gsi = gsi_start_bb (else_bb);
+      stmt = gimple_build_assign (tmp_var, goacc_flags);
+      gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
+
+      make_edge (cond_bb, then_bb, EDGE_TRUE_VALUE);
+      make_edge (cond_bb, else_bb, EDGE_FALSE_VALUE);
+      add_bb_to_loop (then_bb, cond_bb->loop_father);
+      add_bb_to_loop (else_bb, cond_bb->loop_father);
+      make_edge (then_bb, new_bb, EDGE_FALLTHRU);
+      make_edge (else_bb, new_bb, EDGE_FALLTHRU);
+
+      goacc_flags = tmp_var;
+      gsi = gsi_last_nondebug_bb (new_bb);
+    }
+
   if (need_device_adjustment)
     {
       tree uns = fold_convert (unsigned_type_node, device);
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 1857b5b960f..d2adf776939 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -1493,6 +1493,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 
 	case OMP_CLAUSE_FINAL:
 	case OMP_CLAUSE_IF:
+	case OMP_CLAUSE_SELF:
 	case OMP_CLAUSE_NUM_THREADS:
 	case OMP_CLAUSE_NUM_TEAMS:
 	case OMP_CLAUSE_THREAD_LIMIT:
@@ -1920,6 +1921,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	case OMP_CLAUSE_COPYIN:
 	case OMP_CLAUSE_DEFAULT:
 	case OMP_CLAUSE_IF:
+	case OMP_CLAUSE_SELF:
 	case OMP_CLAUSE_NUM_THREADS:
 	case OMP_CLAUSE_NUM_TEAMS:
 	case OMP_CLAUSE_THREAD_LIMIT:
diff --git a/gcc/testsuite/c-c++-common/goacc/self-clause-1.c b/gcc/testsuite/c-c++-common/goacc/self-clause-1.c
new file mode 100644
index 00000000000..ed5d072e81f
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/self-clause-1.c
@@ -0,0 +1,22 @@
+/* { dg-skip-if "not yet" { c++ } } */
+
+void
+f (int b)
+{
+  struct { int i; } *p;
+
+#pragma acc parallel self self(b) /* { dg-error "too many 'self' clauses" } */
+  ;
+#pragma acc parallel self(*p) /* { dg-error "used struct type value where scalar is required" } */
+  ;
+
+#pragma acc kernels self self(b) /* { dg-error "too many 'self' clauses" } */
+  ;
+#pragma acc kernels self(*p) /* { dg-error "used struct type value where scalar is required" } */
+  ;
+
+#pragma acc serial self self(b) /* { dg-error "too many 'self' clauses" } */
+  ;
+#pragma acc serial self(*p) /* { dg-error "used struct type value where scalar is required" } */
+  ;
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/self-clause-2.c b/gcc/testsuite/c-c++-common/goacc/self-clause-2.c
new file mode 100644
index 00000000000..d932ac9a4a6
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/self-clause-2.c
@@ -0,0 +1,17 @@
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+void
+f (short c)
+{
+#pragma acc parallel self(c) copy(c)
+  /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_parallel map\(tofrom:c \[len: [0-9]+\]\) self\(_[0-9]+\)$} 1 "gimple" } } */
+  ++c;
+
+#pragma acc kernels self(c) copy(c)
+  /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_kernels map\(tofrom:c \[len: [0-9]+\]\) self\(_[0-9]+\)$} 1 "gimple" } } */
+  ++c;
+
+#pragma acc serial self(c) copy(c)
+  /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_serial map\(tofrom:c \[len: [0-9]+\]\) self\(_[0-9]+\)$} 1 "gimple" } } */
+  ++c;
+}
diff --git a/gcc/testsuite/gfortran.dg/goacc/self.f95 b/gcc/testsuite/gfortran.dg/goacc/self.f95
new file mode 100644
index 00000000000..4817f16be56
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/self.f95
@@ -0,0 +1,53 @@
+! { dg-do compile } 
+
+program test
+  implicit none
+
+  logical :: x
+  integer :: i
+
+  !$acc parallel self () ! { dg-error "Invalid character" }
+  !$acc parallel self (i) ! { dg-error "scalar LOGICAL expression" }
+  !$acc end parallel 
+  !$acc parallel self (1) ! { dg-error "scalar LOGICAL expression" }
+  !$acc end parallel 
+
+  !$acc kernels self () ! { dg-error "Invalid character" }
+  !$acc kernels self (i) ! { dg-error "scalar LOGICAL expression" }
+  !$acc end kernels 
+  !$acc kernels self (1) ! { dg-error "scalar LOGICAL expression" }
+  !$acc end kernels
+
+  !$acc serial self () ! { dg-error "Invalid character" }
+  !$acc serial self (i) ! { dg-error "scalar LOGICAL expression" }
+  !$acc end serial
+  !$acc serial self (1) ! { dg-error "scalar LOGICAL expression" }
+  !$acc end serial
+  
+  ! at most one self clause may appear
+  !$acc parallel self (.false.) self (.false.) { dg-error "Duplicated 'self' clause" }
+  !$acc kernels self (.false.) self (.false.) { dg-error "Duplicated 'self' clause" }
+  !$acc serial self (.false.) self (.false.) { dg-error "Duplicated 'self' clause" }
+
+  !$acc parallel self (x)
+  !$acc end parallel
+  !$acc parallel self (.true.)
+  !$acc end parallel
+  !$acc parallel self (i.gt.1)
+  !$acc end parallel
+
+  !$acc kernels self (x)
+  !$acc end kernels
+  !$acc kernels self (.true.)
+  !$acc end kernels
+  !$acc kernels self (i.gt.1)
+  !$acc end kernels
+
+  !$acc serial self (x)
+  !$acc end serial
+  !$acc serial self (.true.)
+  !$acc end serial
+  !$acc serial self (i.gt.1)
+  !$acc end serial
+
+end program test
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index c48a12b378f..3fd5a3dff78 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -524,6 +524,9 @@ enum omp_clause_code {
 
   /* OpenACC clause: nohost.  */
   OMP_CLAUSE_NOHOST,
+
+  /* OpenACC clause: self.  */
+  OMP_CLAUSE_SELF,
 };
 
 #undef DEFTREESTRUCT
diff --git a/gcc/tree-nested.cc b/gcc/tree-nested.cc
index ae7d1f1f6a8..8dd4c8acf25 100644
--- a/gcc/tree-nested.cc
+++ b/gcc/tree-nested.cc
@@ -1366,6 +1366,7 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
 	  /* FALLTHRU */
 	case OMP_CLAUSE_FINAL:
 	case OMP_CLAUSE_IF:
+	case OMP_CLAUSE_SELF:
 	case OMP_CLAUSE_NUM_THREADS:
 	case OMP_CLAUSE_DEPEND:
 	case OMP_CLAUSE_DOACROSS:
@@ -2156,6 +2157,7 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
 	  /* FALLTHRU */
 	case OMP_CLAUSE_FINAL:
 	case OMP_CLAUSE_IF:
+	case OMP_CLAUSE_SELF:
 	case OMP_CLAUSE_NUM_THREADS:
 	case OMP_CLAUSE_DEPEND:
 	case OMP_CLAUSE_DOACROSS:
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index 25d191b10fd..a743e3cdfd8 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -1450,7 +1450,12 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
 			 false);
       pp_right_paren (pp);
       break;
-
+    case OMP_CLAUSE_SELF:
+      pp_string (pp, "self(");
+      dump_generic_node (pp, OMP_CLAUSE_SELF_EXPR (clause),
+			 spc, flags, false);
+      pp_right_paren (pp);
+      break;
     default:
       gcc_unreachable ();
     }
diff --git a/gcc/tree.cc b/gcc/tree.cc
index 8e144bc090e..18a27d2faf9 100644
--- a/gcc/tree.cc
+++ b/gcc/tree.cc
@@ -326,6 +326,7 @@ unsigned const char omp_clause_num_ops[] =
   0, /* OMP_CLAUSE_IF_PRESENT */
   0, /* OMP_CLAUSE_FINALIZE */
   0, /* OMP_CLAUSE_NOHOST */
+  1, /* OMP_CLAUSE_SELF */
 };
 
 const char * const omp_clause_code_name[] =
@@ -417,6 +418,7 @@ const char * const omp_clause_code_name[] =
   "if_present",
   "finalize",
   "nohost",
+  "self",
 };
 
 /* Unless specific to OpenACC, we tend to internally maintain OpenMP-centric
diff --git a/gcc/tree.h b/gcc/tree.h
index 1854fe4a7d4..f04457022ac 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1711,6 +1711,8 @@ class auto_suppress_location_wrappers
   OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_HINT), 0)
 #define OMP_CLAUSE_FILTER_EXPR(NODE) \
   OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_FILTER), 0)
+#define OMP_CLAUSE_SELF_EXPR(NODE) \
+  OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_SELF), 0)
 
 #define OMP_CLAUSE_GRAINSIZE_EXPR(NODE) \
   OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_GRAINSIZE),0)
diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index 49b7dd86ff5..8a349ca504f 100644
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -305,6 +305,8 @@ enum gomp_map_kind
 
 /* Force host fallback execution.  */
 #define GOACC_FLAG_HOST_FALLBACK	(1 << 0)
+/* Execute on local device (i.e. host multicore CPU).  */
+#define GOACC_FLAG_LOCAL_DEVICE 	(1 << 1)
 
 /* For legacy reasons, in the ABI, the GOACC_FLAGs are encoded as an inverted
    bitmask.  */
diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
index 363e6656982..cf37a1bdd7d 100644
--- a/libgomp/oacc-parallel.c
+++ b/libgomp/oacc-parallel.c
@@ -193,6 +193,17 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (void *),
       goacc_restore_bind ();
       goto out_prof;
     }
+  else if (flags & GOACC_FLAG_LOCAL_DEVICE)
+    {
+      /* TODO: a proper pthreads based "multi-core CPU" local device
+	 implementation. Currently, this is still the same as host-fallback.  */
+      prof_info.device_type = acc_device_host;
+      api_info.device_type = prof_info.device_type;
+      goacc_save_and_set_bind (acc_device_host);
+      fn (hostaddrs);
+      goacc_restore_bind ();
+      goto out_prof;
+    }
   else if (acc_device_type (acc_dev->type) == acc_device_host)
     {
       fn (hostaddrs);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/self-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/self-1.c
new file mode 100644
index 00000000000..752e16e8545
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/self-1.c
@@ -0,0 +1,962 @@
+#include <openacc.h>
+#include <stdlib.h>
+#include <stdbool.h>
+
+#define N   32
+
+int
+main(int argc, char **argv)
+{
+    float *a, *b, *d_a, *d_b, exp, exp2;
+    int i;
+    const int one = 1;
+    const int zero = 0;
+    int n;
+
+    a = (float *) malloc (N * sizeof (float));
+    b = (float *) malloc (N * sizeof (float));
+    d_a = (float *) acc_malloc (N * sizeof (float));
+    d_b = (float *) acc_malloc (N * sizeof (float));
+
+    for (i = 0; i < N; i++)
+        a[i] = 4.0;
+
+#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) self(0)
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+        {
+            if (acc_on_device (acc_device_host))
+                b[ii] = a[ii] + 1;
+            else
+                b[ii] = a[ii];
+        }
+    }
+
+#if ACC_MEM_SHARED
+    exp = 5.0;
+#else
+    exp = 4.0;
+#endif
+
+    for (i = 0; i < N; i++)
+    {
+        if (b[i] != exp)
+            abort();
+    }
+
+    for (i = 0; i < N; i++)
+        a[i] = 16.0;
+
+#pragma acc parallel self(1)
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+        {
+            if (acc_on_device (acc_device_host))
+                b[ii] = a[ii] + 1;
+            else
+                b[ii] = a[ii];
+        }
+    }
+
+    for (i = 0; i < N; i++)
+    {
+        if (b[i] != 17.0)
+            abort();
+    }
+
+    for (i = 0; i < N; i++)
+        a[i] = 8.0;
+
+#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) self(!one)
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+        {
+            if (acc_on_device (acc_device_host))
+                b[ii] = a[ii] + 1;
+            else
+                b[ii] = a[ii];
+        }
+    }
+
+#if ACC_MEM_SHARED
+    exp = 9.0;
+#else
+    exp = 8.0;
+#endif
+
+    for (i = 0; i < N; i++)
+    {
+        if (b[i] != exp)
+            abort();
+    }
+
+    for (i = 0; i < N; i++)
+        a[i] = 22.0;
+
+#pragma acc parallel self(!zero)
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+        {
+            if (acc_on_device (acc_device_host))
+                b[ii] = a[ii] + 1;
+            else
+                b[ii] = a[ii];
+        }
+    }
+
+    for (i = 0; i < N; i++)
+    {
+        if (b[i] != 23.0)
+            abort();
+    }
+
+    for (i = 0; i < N; i++)
+        a[i] = 16.0;
+
+#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) self(false)
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+        {
+            if (acc_on_device (acc_device_host))
+                b[ii] = a[ii] + 1;
+            else
+                b[ii] = a[ii];
+        }
+    }
+
+#if ACC_MEM_SHARED
+    exp = 17.0;
+#else
+    exp = 16.0;
+#endif
+
+    for (i = 0; i < N; i++)
+    {
+        if (b[i] != exp)
+            abort();
+    }
+
+    for (i = 0; i < N; i++)
+        a[i] = 76.0;
+
+#pragma acc parallel self(true)
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+        {
+            if (acc_on_device (acc_device_host))
+                b[ii] = a[ii] + 1;
+            else
+                b[ii] = a[ii];
+        }
+    }
+
+    for (i = 0; i < N; i++)
+    {
+        if (b[i] != 77.0)
+            abort();
+    }
+
+    for (i = 0; i < N; i++)
+        a[i] = 22.0;
+
+    n = 1;
+
+#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) self(!n)
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+        {
+            if (acc_on_device (acc_device_host))
+                b[ii] = a[ii] + 1;
+            else
+                b[ii] = a[ii];
+        }
+    }
+
+#if ACC_MEM_SHARED
+    exp = 23.0;
+#else
+    exp = 22.0;
+#endif
+
+    for (i = 0; i < N; i++)
+    {
+        if (b[i] != exp)
+            abort();
+    }
+
+    for (i = 0; i < N; i++)
+        a[i] = 18.0;
+
+    n = 0;
+
+#pragma acc parallel self(!n)
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+        {
+            if (acc_on_device (acc_device_host))
+                b[ii] = a[ii] + 1;
+            else
+                b[ii] = a[ii];
+        }
+    }
+
+    for (i = 0; i < N; i++)
+    {
+        if (b[i] != 19.0)
+            abort();
+    }
+
+    for (i = 0; i < N; i++)
+        a[i] = 49.0;
+
+    n = 1;
+
+#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) self(!(n + n))
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+        {
+            if (acc_on_device (acc_device_host))
+                b[ii] = a[ii] + 1;
+            else
+                b[ii] = a[ii];
+        }
+    }
+
+#if ACC_MEM_SHARED
+    exp = 50.0;
+#else
+    exp = 49.0;
+#endif
+
+    for (i = 0; i < N; i++)
+    {
+        if (b[i] != exp)
+            abort();
+    }
+
+    for (i = 0; i < N; i++)
+        a[i] = 38.0;
+
+    n = 0;
+
+#pragma acc parallel self(!(n + n))
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+        {
+            if (acc_on_device (acc_device_host))
+                b[ii] = a[ii] + 1;
+            else
+                b[ii] = a[ii];
+        }
+    }
+
+    for (i = 0; i < N; i++)
+    {
+        if (b[i] != 39.0)
+            abort();
+    }
+
+    for (i = 0; i < N; i++)
+        a[i] = 91.0;
+
+#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) self(!(-2))
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+        {
+            if (acc_on_device (acc_device_host))
+                b[ii] = a[ii] + 1;
+            else
+                b[ii] = a[ii];
+        }
+    }
+
+#if ACC_MEM_SHARED
+    exp = 92.0;
+#else
+    exp = 91.0;
+#endif
+
+    for (i = 0; i < N; i++)
+    {
+        if (b[i] != exp)
+            abort();
+    }
+
+    for (i = 0; i < N; i++)
+        a[i] = 43.0;
+
+#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) self(one != 1)
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+        {
+            if (acc_on_device (acc_device_host))
+                b[ii] = a[ii] + 1;
+            else
+                b[ii] = a[ii];
+        }
+    }
+
+#if ACC_MEM_SHARED
+    exp = 44.0;
+#else
+    exp = 43.0;
+#endif
+
+    for (i = 0; i < N; i++)
+    {
+        if (b[i] != exp)
+            abort();
+    }
+
+    for (i = 0; i < N; i++)
+        a[i] = 87.0;
+
+#pragma acc parallel self(one != 0)
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+        {
+            if (acc_on_device (acc_device_host))
+                b[ii] = a[ii] + 1;
+            else
+                b[ii] = a[ii];
+        }
+    }
+
+    for (i = 0; i < N; i++)
+    {
+        if (b[i] != 88.0)
+            abort();
+    }
+
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 3.0;
+        b[i] = 9.0;
+    }
+
+#if ACC_MEM_SHARED
+    exp = 0.0;
+    exp2 = 0.0;
+#else
+    acc_map_data (a, d_a, N * sizeof (float));
+    acc_map_data (b, d_b, N * sizeof (float));
+    exp = 3.0;
+    exp2 = 9.0;
+#endif
+
+#pragma acc update device(a[0:N], b[0:N]) if(1)
+
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 0.0;
+        b[i] = 0.0;
+    }
+
+#pragma acc update host(a[0:N], b[0:N]) if(1)
+
+    for (i = 0; i < N; i++)
+    {
+        if (a[i] != exp)
+            abort();
+
+        if (b[i] != exp2)
+            abort();
+    }
+
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 6.0;
+        b[i] = 12.0;
+    }
+
+#pragma acc update device(a[0:N], b[0:N]) if(0)
+
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 0.0;
+        b[i] = 0.0;
+    }
+
+#pragma acc update host(a[0:N], b[0:N]) if(1)
+
+    for (i = 0; i < N; i++)
+    {
+        if (a[i] != exp)
+            abort();
+
+        if (b[i] != exp2)
+            abort();
+    }
+
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 26.0;
+        b[i] = 21.0;
+    }
+
+#pragma acc update device(a[0:N], b[0:N]) if(1)
+
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 0.0;
+        b[i] = 0.0;
+    }
+
+#pragma acc update host(a[0:N], b[0:N]) if(0)
+
+    for (i = 0; i < N; i++)
+    {
+        if (a[i] != 0.0)
+            abort();
+
+        if (b[i] != 0.0)
+            abort();
+    }
+
+#if !ACC_MEM_SHARED
+    acc_unmap_data (a);
+    acc_unmap_data (b);
+#endif
+
+    acc_free (d_a);
+    acc_free (d_b);
+
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 4.0;
+        b[i] = 0.0;
+    }
+
+#pragma acc data copyin(a[0:N]) copyout(b[0:N]) if(1)
+{
+#pragma acc parallel present(a[0:N])
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+        {
+            b[ii] = a[ii];
+        }
+    }
+}
+
+    for (i = 0; i < N; i++)
+    {
+        if (b[i] != 4.0)
+            abort();
+    }
+
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 8.0;
+        b[i] = 1.0;
+    }
+
+#pragma acc data copyin(a[0:N]) copyout(b[0:N]) if(0)
+{
+#if !ACC_MEM_SHARED
+    if (acc_is_present (a, N * sizeof (float)))
+        abort ();
+#endif
+
+#if !ACC_MEM_SHARED
+    if (acc_is_present (b, N * sizeof (float)))
+        abort ();
+#endif
+}
+
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 18.0;
+        b[i] = 21.0;
+    }
+
+#pragma acc data copyin(a[0:N]) if(1)
+{
+#if !ACC_MEM_SHARED
+    if (!acc_is_present (a, N * sizeof (float)))
+        abort ();
+#endif
+
+#pragma acc data copyout(b[0:N]) if(0)
+    {
+#if !ACC_MEM_SHARED
+        if (acc_is_present (b, N * sizeof (float)))
+            abort ();
+#endif
+
+#pragma acc data copyout(b[0:N]) if(1)
+        {
+#pragma acc parallel present(a[0:N]) present(b[0:N])
+            {
+                int ii;
+
+                for (ii = 0; ii < N; ii++)
+                {
+                    b[ii] = a[ii];
+                }
+            }
+        }
+
+#if !ACC_MEM_SHARED
+        if (acc_is_present (b, N * sizeof (float)))
+            abort ();
+#endif
+    }
+}
+
+    for (i = 0; i < N; i++)
+    {
+        if (b[i] != 18.0)
+            abort ();
+    }
+
+#pragma acc enter data copyin (b[0:N]) if (0)
+
+#if !ACC_MEM_SHARED
+    if (acc_is_present (b, N * sizeof (float)))
+	abort ();
+#endif
+
+#pragma acc exit data delete (b[0:N]) if (0)
+
+#pragma acc enter data copyin (b[0:N]) if (1)
+
+#if !ACC_MEM_SHARED
+    if (!acc_is_present (b, N * sizeof (float)))
+	abort ();
+#endif
+
+#pragma acc exit data delete (b[0:N]) if (1)
+
+#if !ACC_MEM_SHARED
+    if (acc_is_present (b, N * sizeof (float)))
+	abort ();
+#endif
+
+#pragma acc enter data copyin (b[0:N]) if (zero)
+
+#if !ACC_MEM_SHARED
+    if (acc_is_present (b, N * sizeof (float)))
+	abort ();
+#endif
+
+#pragma acc exit data delete (b[0:N]) if (zero)
+
+#pragma acc enter data copyin (b[0:N]) if (one)
+
+#if !ACC_MEM_SHARED
+    if (!acc_is_present (b, N * sizeof (float)))
+	abort ();
+#endif
+
+#pragma acc exit data delete (b[0:N]) if (one)
+
+#if !ACC_MEM_SHARED
+    if (acc_is_present (b, N * sizeof (float)))
+	abort ();
+#endif
+
+#pragma acc enter data copyin (b[0:N]) if (one == 0)
+
+#if !ACC_MEM_SHARED
+    if (acc_is_present (b, N * sizeof (float)))
+	abort ();
+#endif
+
+#pragma acc exit data delete (b[0:N]) if (one == 0)
+
+#pragma acc enter data copyin (b[0:N]) if (one == 1)
+
+#if !ACC_MEM_SHARED
+    if (!acc_is_present (b, N * sizeof (float)))
+	abort ();
+#endif
+
+#pragma acc exit data delete (b[0:N]) if (one == 1)
+
+#if !ACC_MEM_SHARED
+    if (acc_is_present (b, N * sizeof (float)))
+	abort ();
+#endif
+
+    for (i = 0; i < N; i++)
+        a[i] = 4.0;
+
+#pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) self(0)
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+        {
+            if (acc_on_device (acc_device_host))
+                b[ii] = a[ii] + 1;
+            else
+                b[ii] = a[ii];
+        }
+    }
+
+#if ACC_MEM_SHARED
+    exp = 5.0;
+#else
+    exp = 4.0;
+#endif
+
+    for (i = 0; i < N; i++)
+    {
+        if (b[i] != exp)
+            abort();
+    }
+
+    for (i = 0; i < N; i++)
+        a[i] = 16.0;
+
+#pragma acc kernels self(1)
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+        {
+            if (acc_on_device (acc_device_host))
+                b[ii] = a[ii] + 1;
+            else
+                b[ii] = a[ii];
+        }
+    }
+
+    for (i = 0; i < N; i++)
+    {
+        if (b[i] != 17.0)
+            abort();
+    }
+
+    for (i = 0; i < N; i++)
+        a[i] = 8.0;
+
+#pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) self(!one)
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+        {
+            if (acc_on_device (acc_device_host))
+                b[ii] = a[ii] + 1;
+            else
+                b[ii] = a[ii];
+        }
+    }
+
+#if ACC_MEM_SHARED
+    exp = 9.0;
+#else
+    exp = 8.0;
+#endif
+
+    for (i = 0; i < N; i++)
+    {
+        if (b[i] != exp)
+            abort();
+    }
+
+    for (i = 0; i < N; i++)
+        a[i] = 22.0;
+
+#pragma acc kernels self(!zero)
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+        {
+            if (acc_on_device (acc_device_host))
+                b[ii] = a[ii] + 1;
+            else
+                b[ii] = a[ii];
+        }
+    }
+
+    for (i = 0; i < N; i++)
+    {
+        if (b[i] != 23.0)
+            abort();
+    }
+
+    for (i = 0; i < N; i++)
+        a[i] = 16.0;
+
+#pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) self(false)
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+        {
+            if (acc_on_device (acc_device_host))
+                b[ii] = a[ii] + 1;
+            else
+                b[ii] = a[ii];
+        }
+    }
+
+#if ACC_MEM_SHARED
+    exp = 17.0;
+#else
+    exp = 16.0;
+#endif
+
+    for (i = 0; i < N; i++)
+    {
+        if (b[i] != exp)
+            abort();
+    }
+
+    for (i = 0; i < N; i++)
+        a[i] = 76.0;
+
+#pragma acc kernels self(true)
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+        {
+            if (acc_on_device (acc_device_host))
+                b[ii] = a[ii] + 1;
+            else
+                b[ii] = a[ii];
+        }
+    }
+
+    for (i = 0; i < N; i++)
+    {
+        if (b[i] != 77.0)
+            abort();
+    }
+
+    for (i = 0; i < N; i++)
+        a[i] = 22.0;
+
+    n = 1;
+
+#pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) self(!n)
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+        {
+            if (acc_on_device (acc_device_host))
+                b[ii] = a[ii] + 1;
+            else
+                b[ii] = a[ii];
+        }
+    }
+
+#if ACC_MEM_SHARED
+    exp = 23.0;
+#else
+    exp = 22.0;
+#endif
+
+    for (i = 0; i < N; i++)
+    {
+        if (b[i] != exp)
+            abort();
+    }
+
+    for (i = 0; i < N; i++)
+        a[i] = 18.0;
+
+    n = 0;
+
+#pragma acc kernels self(!n)
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+        {
+            if (acc_on_device (acc_device_host))
+                b[ii] = a[ii] + 1;
+            else
+                b[ii] = a[ii];
+        }
+    }
+
+    for (i = 0; i < N; i++)
+    {
+        if (b[i] != 19.0)
+            abort();
+    }
+
+    for (i = 0; i < N; i++)
+        a[i] = 49.0;
+
+    n = 1;
+
+#pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) self((n + n) == 0)
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+        {
+            if (acc_on_device (acc_device_host))
+                b[ii] = a[ii] + 1;
+            else
+                b[ii] = a[ii];
+        }
+    }
+
+#if ACC_MEM_SHARED
+    exp = 50.0;
+#else
+    exp = 49.0;
+#endif
+
+    for (i = 0; i < N; i++)
+    {
+        if (b[i] != exp)
+            abort();
+    }
+
+    for (i = 0; i < N; i++)
+        a[i] = 38.0;
+
+    n = 0;
+
+#pragma acc kernels self(!(n + n))
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+        {
+            if (acc_on_device (acc_device_host))
+                b[ii] = a[ii] + 1;
+            else
+                b[ii] = a[ii];
+        }
+    }
+
+    for (i = 0; i < N; i++)
+    {
+        if (b[i] != 39.0)
+            abort();
+    }
+
+    for (i = 0; i < N; i++)
+        a[i] = 91.0;
+
+#pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) self(!(-2))
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+        {
+            if (acc_on_device (acc_device_host))
+                b[ii] = a[ii] + 1;
+            else
+                b[ii] = a[ii];
+        }
+    }
+
+#if ACC_MEM_SHARED
+    exp = 92.0;
+#else
+    exp = 91.0;
+#endif
+
+    for (i = 0; i < N; i++)
+    {
+        if (b[i] != exp)
+            abort();
+    }
+
+    for (i = 0; i < N; i++)
+        a[i] = 43.0;
+
+#pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) self(one != 1)
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+        {
+            if (acc_on_device (acc_device_host))
+                b[ii] = a[ii] + 1;
+            else
+                b[ii] = a[ii];
+        }
+    }
+
+#if ACC_MEM_SHARED
+    exp = 44.0;
+#else
+    exp = 43.0;
+#endif
+
+    for (i = 0; i < N; i++)
+    {
+        if (b[i] != exp)
+            abort();
+    }
+
+    for (i = 0; i < N; i++)
+        a[i] = 87.0;
+
+#pragma acc kernels self(one != 0)
+    {
+        int ii;
+
+        for (ii = 0; ii < N; ii++)
+        {
+            if (acc_on_device (acc_device_host))
+                b[ii] = a[ii] + 1;
+            else
+                b[ii] = a[ii];
+        }
+    }
+
+    for (i = 0; i < N; i++)
+    {
+        if (b[i] != 88.0)
+            abort();
+    }
+
+    for (i = 0; i < N; i++)
+    {
+        a[i] = 3.0;
+        b[i] = 9.0;
+    }
+
+#if ACC_MEM_SHARED
+    exp = 0.0;
+    exp2 = 0.0;
+#else
+    acc_map_data (a, d_a, N * sizeof (float));
+    acc_map_data (b, d_b, N * sizeof (float));
+    exp = 3.0;
+    exp2 = 9.0;
+#endif
+
+    return 0;
+}
