OpenMP: Add ME support for 'omp allocate' stack variables

Message ID 98cc1611-7369-4a2b-a7df-73200eded3c9@codesourcery.com
State New
Headers
Series OpenMP: Add ME support for 'omp allocate' stack variables |

Checks

Context Check Description
linaro-tcwg-bot/tcwg_gcc_check--master-arm success Testing passed
linaro-tcwg-bot/tcwg_gcc_check--master-aarch64 success Testing passed
linaro-tcwg-bot/tcwg_gcc_build--master-arm success Testing passed
linaro-tcwg-bot/tcwg_gcc_build--master-aarch64 success Testing passed

Commit Message

Tobias Burnus Sept. 18, 2023, 12:22 p.m. UTC
  The attached patch now actually adds GOMP_alloc/free calls for 'omp allocate'.

Besides the addition of the calls and the value expression, it also had to deal with
(implicit) mapping/privatization - such that 'default(none)' did not wrongly trigger
for the value expression (and categorizes the vars correctly for default/defaultmap)
and that mapping/privatization is handled correctly.

Build and regtested on x86-64-gnu-linux (w/o offloading configured, but I tested
separately the libgomp.c/allocate-*.c with nvptx offloading).

Comments, suggestions, remarks?

Tobias

PS: I am sure that some cases are not correctly handled, the question is only which
ones. For C++, I surely should check the cases handled by cxx_omp_disregard_value_expr.
For Fortran, I have to check what should be handled in the FE and what in the ME; in
any case, gfc_omp_disregard_value_expr also exists.
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
  

Comments

Jakub Jelinek Sept. 20, 2023, 1:08 p.m. UTC | #1
On Mon, Sep 18, 2023 at 02:22:50PM +0200, Tobias Burnus wrote:
> The attached patch now actually adds GOMP_alloc/free calls for 'omp allocate'.
> 
> Besides the addition of the calls and the value expression, it also had to deal with
> (implicit) mapping/privatization - such that 'default(none)' did not wrongly trigger
> for the value expression (and categorizes the vars correctly for default/defaultmap)
> and that mapping/privatization is handled correctly.
> 
> Build and regtested on x86-64-gnu-linux (w/o offloading configured, but I tested
> separately the libgomp.c/allocate-*.c with nvptx offloading).
> 
> Comments, suggestions, remarks?

LGTM.

	Jakub
  
Thomas Schwinge Oct. 18, 2023, 9:44 a.m. UTC | #2
Hi Tobias!

No need to change anything now, but in case that's useful later:

On 2023-09-18T14:22:50+0200, Tobias Burnus <tobias@codesourcery.com> wrote:
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.c/allocate-4.c
> @@ -0,0 +1,84 @@
> +/* TODO: move to ../libgomp.c-c++-common once C++ is implemented. */
> +/* NOTE: { target c } is unsupported with with the C compiler.  */
> +[...]

..., just noting that '{ target c }', '{ target c++ }' are trivial to
implement; see libgomp OpenACC testing:

libgomp/testsuite/libgomp.oacc-c/c.exp:proc check_effective_target_c { } {
libgomp/testsuite/libgomp.oacc-c/c.exp-    return 1
libgomp/testsuite/libgomp.oacc-c/c.exp-}
libgomp/testsuite/libgomp.oacc-c/c.exp:proc check_effective_target_c++ { } {
libgomp/testsuite/libgomp.oacc-c/c.exp-    return 0
libgomp/testsuite/libgomp.oacc-c/c.exp-}

libgomp/testsuite/libgomp.oacc-c++/c++.exp:proc check_effective_target_c { } {
libgomp/testsuite/libgomp.oacc-c++/c++.exp-    return 0
libgomp/testsuite/libgomp.oacc-c++/c++.exp-}
libgomp/testsuite/libgomp.oacc-c++/c++.exp:proc check_effective_target_c++ { } {
libgomp/testsuite/libgomp.oacc-c++/c++.exp-    return 1
libgomp/testsuite/libgomp.oacc-c++/c++.exp-}


Grüße
 Thomas


> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.c/allocate-5.c
> @@ -0,0 +1,126 @@
> +/* TODO: move to ../libgomp.c-c++-common once C++ is implemented. */
> +/* NOTE: { target c } is unsupported with with the C compiler.  */
> +[...]

> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.c/allocate-6.c
> @@ -0,0 +1,319 @@
> +/* TODO: move to ../libgomp.c-c++-common once C++ is implemented. */
> +/* NOTE: { target c } is unsupported with with the C compiler.  */
> +[...]
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
  
Tobias Burnus Oct. 18, 2023, 9:53 a.m. UTC | #3
Hi Thomas,

On 18.10.23 11:44, Thomas Schwinge wrote:
> No need to change anything now, but in case that's useful later:
> [...]
> ..., just noting that '{ target c }', '{ target c++ }' are trivial to
> implement; see libgomp OpenACC testing:
>
> libgomp/testsuite/libgomp.oacc-c/c.exp:proc check_effective_target_c { } {

Shouldn't we rather do something like the following (untested)?

--- a/libgomp/testsuite/lib/libgomp.exp
+++ b/libgomp/testsuite/lib/libgomp.exp
@@ -26,2 +26,3 @@ load_gcc_lib timeout.exp
  load_gcc_lib file-format.exp
+load_gcc_lib target-supports.exp
  load_gcc_lib target-supports-dg.expTobias

-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
  
Thomas Schwinge Oct. 18, 2023, 10:01 a.m. UTC | #4
Hi Tobias!

On 2023-10-18T11:53:30+0200, Tobias Burnus <tobias@codesourcery.com> wrote:
> On 18.10.23 11:44, Thomas Schwinge wrote:
>> No need to change anything now, but in case that's useful later:
>> [...]
>> ..., just noting that '{ target c }', '{ target c++ }' are trivial to
>> implement; see libgomp OpenACC testing:
>>
>> libgomp/testsuite/libgomp.oacc-c/c.exp:proc check_effective_target_c { } {
>
> Shouldn't we rather do something like the following (untested)?
>
> --- a/libgomp/testsuite/lib/libgomp.exp
> +++ b/libgomp/testsuite/lib/libgomp.exp
> @@ -26,2 +26,3 @@ load_gcc_lib timeout.exp
>   load_gcc_lib file-format.exp
> +load_gcc_lib target-supports.exp
>   load_gcc_lib target-supports-dg.exp

'gcc/testsuite/lib/target-supports.exp' defines:

    # Return 1 if the language for the compiler under test is C.

    proc check_effective_target_c { } {
        global tool
        if [string match $tool "gcc"] {
            return 1
        }
     return 0
    }

    # Return 1 if the language for the compiler under test is C++.

    proc check_effective_target_c++ { } {
        global tool
        if { [string match $tool "g++"] || [string match $tool "libstdc++"] } {
            return 1
        }
     return 0
    }

However, (per my understanding; not verified) 'tool == libgomp' for
libgomp testing, so that doesn't work.


Grüße
 Thomas
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
  

Patch

OpenMP: Add ME support for 'omp allocate' stack variables

Call GOMP_alloc/free for 'omp allocate' allocated variables. This is
for C only as C++ and Fortran show a sorry already in the FE. Note that
this only applies to stack variables as the C FE shows a sorry for
static variables.

gcc/ChangeLog:

	* gimplify.cc (gimplify_bind_expr): Call GOMP_alloc/free for
	'omp allocate' variables; move stack cleanup after other
	cleanup.
	(omp_notice_variable): Process original decl when decl
	of the value-expression for a 'omp allocate' variable is passed.
	* omp-low.cc (scan_omp_1_op): Handle 'omp allocate' variables

libgomp/ChangeLog:

	* libgomp.texi (OpenMP 5.1 Impl.): Mark 'omp allocate' as
	implemented for C only.
	* testsuite/libgomp.c/allocate-4.c: New test.
	* testsuite/libgomp.c/allocate-5.c: New test.
	* testsuite/libgomp.c/allocate-6.c: New test.

gcc/testsuite/ChangeLog:

	* c-c++-common/gomp/allocate-11.c: Remove C-only dg-message
	for 'sorry, unimplemented'.
	* c-c++-common/gomp/allocate-12.c: Likewise.
	* c-c++-common/gomp/allocate-15.c: Likewise.
	* c-c++-common/gomp/allocate-9.c: Likewise.
	* c-c++-common/gomp/allocate-10.c: New test.
	* c-c++-common/gomp/allocate-17.c: New test.

 gcc/gimplify.cc                               | 108 +++++++--
 gcc/omp-low.cc                                |  28 ++-
 gcc/testsuite/c-c++-common/gomp/allocate-10.c |  49 ++++
 gcc/testsuite/c-c++-common/gomp/allocate-11.c |   3 -
 gcc/testsuite/c-c++-common/gomp/allocate-12.c |   3 -
 gcc/testsuite/c-c++-common/gomp/allocate-15.c |   4 +-
 gcc/testsuite/c-c++-common/gomp/allocate-17.c |  37 +++
 gcc/testsuite/c-c++-common/gomp/allocate-9.c  |   2 -
 libgomp/libgomp.texi                          |   5 +-
 libgomp/testsuite/libgomp.c/allocate-4.c      |  84 +++++++
 libgomp/testsuite/libgomp.c/allocate-5.c      | 126 ++++++++++
 libgomp/testsuite/libgomp.c/allocate-6.c      | 319 ++++++++++++++++++++++++++
 12 files changed, 733 insertions(+), 35 deletions(-)

diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index a0e8cc2199d..9f4722f7458 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -36,6 +36,7 @@  along with GCC; see the file COPYING3.  If not see
 #include "cgraph.h"
 #include "tree-pretty-print.h"
 #include "diagnostic-core.h"
+#include "diagnostic.h"		/* For errorcount.  */
 #include "alias.h"
 #include "fold-const.h"
 #include "calls.h"
@@ -1372,6 +1373,7 @@  gimplify_bind_expr (tree *expr_p, gimple_seq *pre_p)
 	      && (attr = lookup_attribute ("omp allocate", DECL_ATTRIBUTES (t)))
 		 != NULL_TREE)
 	    {
+	      gcc_assert (!DECL_HAS_VALUE_EXPR_P (t));
 	      tree alloc = TREE_PURPOSE (TREE_VALUE (attr));
 	      tree align = TREE_VALUE (TREE_VALUE (attr));
 	      /* Allocate directives that appear in a target region must specify
@@ -1396,12 +1398,56 @@  gimplify_bind_expr (tree *expr_p, gimple_seq *pre_p)
 		error_at (DECL_SOURCE_LOCATION (t),
 			  "%<allocate%> directive for %qD inside a target "
 			  "region must specify an %<allocator%> clause", t);
-	      else if (align != NULL_TREE
-		       || alloc == NULL_TREE
-		       || !integer_onep (alloc))
-	        sorry_at (DECL_SOURCE_LOCATION (t),
-			  "OpenMP %<allocate%> directive, used for %qD, not "
-			  "yet supported", t);
+	      /* Skip for omp_default_mem_alloc (= 1),
+		 unless align is present. */
+	      else if (!errorcount
+		       && (align != NULL_TREE
+			   || alloc == NULL_TREE
+			   || !integer_onep (alloc)))
+		{
+		  tree tmp = build_pointer_type (TREE_TYPE (t));
+		  tree v = create_tmp_var (tmp, get_name (t));
+		  DECL_IGNORED_P (v) = 0;
+		  tmp = remove_attribute ("omp allocate", DECL_ATTRIBUTES (t));
+		  DECL_ATTRIBUTES (v)
+		    = tree_cons (get_identifier ("omp allocate var"),
+				 build_tree_list (NULL_TREE, t), tmp);
+		  tmp = build_fold_indirect_ref (v);
+		  TREE_THIS_NOTRAP (tmp) = 1;
+		  SET_DECL_VALUE_EXPR (t, tmp);
+		  DECL_HAS_VALUE_EXPR_P (t) = 1;
+		  tree sz = TYPE_SIZE_UNIT (TREE_TYPE (t));
+		  if (alloc == NULL_TREE)
+		    alloc = build_zero_cst (ptr_type_node);
+		  if (align == NULL_TREE)
+		    align = build_int_cst (size_type_node, DECL_ALIGN_UNIT (t));
+		  else
+		    align = build_int_cst (size_type_node,
+					   MAX (tree_to_uhwi (align),
+						DECL_ALIGN_UNIT (t)));
+		  tmp = builtin_decl_explicit (BUILT_IN_GOMP_ALLOC);
+		  tmp = build_call_expr_loc (DECL_SOURCE_LOCATION (t), tmp,
+					     3, align, sz, alloc);
+		  tmp = fold_build2_loc (DECL_SOURCE_LOCATION (t), MODIFY_EXPR,
+					 TREE_TYPE (v), v,
+					 fold_convert (TREE_TYPE (v), tmp));
+		  gcc_assert (BIND_EXPR_BODY (bind_expr) != NULL_TREE
+			      && (TREE_CODE (BIND_EXPR_BODY (bind_expr))
+				  == STATEMENT_LIST));
+		  tree_stmt_iterator e = tsi_start (BIND_EXPR_BODY (bind_expr));
+		  while (!tsi_end_p (e))
+		    {
+		      if ((TREE_CODE (*e) == DECL_EXPR
+			   && TREE_OPERAND (*e, 0) == t)
+			  || (TREE_CODE (*e) == CLEANUP_POINT_EXPR
+			      && TREE_CODE (TREE_OPERAND (*e, 0)) == DECL_EXPR
+			      && TREE_OPERAND (TREE_OPERAND (*e, 0), 0) == t))
+		      break;
+		      ++e;
+		    }
+		  gcc_assert (!tsi_end_p (e));
+		  tsi_link_before (&e, tmp, TSI_SAME_STMT);
+		}
 	    }
 
 	  /* Mark variable as local.  */
@@ -1486,22 +1532,6 @@  gimplify_bind_expr (tree *expr_p, gimple_seq *pre_p)
   cleanup = NULL;
   stack_save = NULL;
 
-  /* If the code both contains VLAs and calls alloca, then we cannot reclaim
-     the stack space allocated to the VLAs.  */
-  if (gimplify_ctxp->save_stack && !gimplify_ctxp->keep_stack)
-    {
-      gcall *stack_restore;
-
-      /* Save stack on entry and restore it on exit.  Add a try_finally
-	 block to achieve this.  */
-      build_stack_save_restore (&stack_save, &stack_restore);
-
-      gimple_set_location (stack_save, start_locus);
-      gimple_set_location (stack_restore, end_locus);
-
-      gimplify_seq_add_stmt (&cleanup, stack_restore);
-    }
-
   /* Add clobbers for all variables that go out of scope.  */
   for (t = BIND_EXPR_VARS (bind_expr); t ; t = DECL_CHAIN (t))
     {
@@ -1509,6 +1539,17 @@  gimplify_bind_expr (tree *expr_p, gimple_seq *pre_p)
 	  && !is_global_var (t)
 	  && DECL_CONTEXT (t) == current_function_decl)
 	{
+	  if (flag_openmp
+	      && DECL_HAS_VALUE_EXPR_P (t)
+	      && TREE_USED (t)
+	      && lookup_attribute ("omp allocate", DECL_ATTRIBUTES (t)))
+	    {
+	      tree tmp = builtin_decl_explicit (BUILT_IN_GOMP_FREE);
+	      tmp = build_call_expr_loc (end_locus, tmp, 2,
+					 TREE_OPERAND (DECL_VALUE_EXPR (t), 0),
+					 build_zero_cst (ptr_type_node));
+	      gimplify_and_add (tmp, &cleanup);
+	    }
 	  if (!DECL_HARD_REGISTER (t)
 	      && !TREE_THIS_VOLATILE (t)
 	      && !DECL_HAS_VALUE_EXPR_P (t)
@@ -1565,6 +1606,22 @@  gimplify_bind_expr (tree *expr_p, gimple_seq *pre_p)
 	gimplify_ctxp->live_switch_vars->remove (t);
     }
 
+  /* If the code both contains VLAs and calls alloca, then we cannot reclaim
+     the stack space allocated to the VLAs.  */
+  if (gimplify_ctxp->save_stack && !gimplify_ctxp->keep_stack)
+    {
+      gcall *stack_restore;
+
+      /* Save stack on entry and restore it on exit.  Add a try_finally
+	 block to achieve this.  */
+      build_stack_save_restore (&stack_save, &stack_restore);
+
+      gimple_set_location (stack_save, start_locus);
+      gimple_set_location (stack_restore, end_locus);
+
+      gimplify_seq_add_stmt (&cleanup, stack_restore);
+    }
+
   if (ret_clauses)
     {
       gomp_target *stmt;
@@ -7894,6 +7951,13 @@  omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
   if (error_operand_p (decl))
     return false;
 
+  if (DECL_ARTIFICIAL (decl))
+    {
+      tree attr = lookup_attribute ("omp allocate var", DECL_ATTRIBUTES (decl));
+      if (attr)
+	decl = TREE_VALUE (TREE_VALUE (attr));
+    }
+
   if (ctx->region_type == ORT_NONE)
     return lang_hooks.decls.omp_disregard_value_expr (decl, false);
 
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 5d7c32dac39..b0c3ef7a9cc 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -3951,6 +3951,7 @@  scan_omp_1_op (tree *tp, int *walk_subtrees, void *data)
   struct walk_stmt_info *wi = (struct walk_stmt_info *) data;
   omp_context *ctx = (omp_context *) wi->info;
   tree t = *tp;
+  tree tmp;
 
   switch (TREE_CODE (t))
     {
@@ -3960,12 +3961,37 @@  scan_omp_1_op (tree *tp, int *walk_subtrees, void *data)
     case RESULT_DECL:
       if (ctx)
 	{
+	  tmp = NULL_TREE;
+	  if (TREE_CODE (t) == VAR_DECL
+	      && (tmp = lookup_attribute ("omp allocate var",
+					  DECL_ATTRIBUTES (t))) != NULL_TREE)
+	    t = TREE_VALUE (TREE_VALUE (tmp));
 	  tree repl = remap_decl (t, &ctx->cb);
 	  gcc_checking_assert (TREE_CODE (repl) != ERROR_MARK);
-	  *tp = repl;
+	  if (tmp != NULL_TREE  && t != repl)
+	    *tp = build_fold_addr_expr (repl);
+	  else if (tmp == NULL_TREE)
+	    *tp = repl;
 	}
       break;
 
+    case INDIRECT_REF:
+    case MEM_REF:
+      if (ctx
+	  && TREE_CODE (TREE_OPERAND (t, 0)) == VAR_DECL
+	  && ((tmp = lookup_attribute ("omp allocate var",
+				       DECL_ATTRIBUTES (TREE_OPERAND (t, 0))))
+	       != NULL_TREE))
+	{
+	  tmp = TREE_VALUE (TREE_VALUE (tmp));
+	  tree repl = remap_decl (tmp, &ctx->cb);
+	  gcc_checking_assert (TREE_CODE (repl) != ERROR_MARK);
+	  if (tmp != repl)
+	    *tp = repl;
+	  break;
+	}
+      gcc_fallthrough ();
+
     default:
       if (ctx && TYPE_P (t))
 	*tp = remap_type (t, &ctx->cb);
diff --git a/gcc/testsuite/c-c++-common/gomp/allocate-10.c b/gcc/testsuite/c-c++-common/gomp/allocate-10.c
new file mode 100644
index 00000000000..7e8f579871c
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/allocate-10.c
@@ -0,0 +1,49 @@ 
+/* TODO: enable for C++ once implemented. */
+/* { dg-do compile { target c } } */
+/* { dg-additional-options "-Wall -fdump-tree-gimple" } */
+
+typedef enum omp_allocator_handle_t
+#if __cplusplus >= 201103L
+: __UINTPTR_TYPE__
+#endif
+{
+  omp_default_mem_alloc = 1,
+  __omp_allocator_handle_t_max__ = __UINTPTR_MAX__
+} omp_allocator_handle_t;
+
+void
+f()
+{
+  int n;
+  int A[n]; /* { dg-warning "'n' is used uninitialized" } */
+  /* { dg-warning "unused variable 'A'" "" { target *-*-* } .-1 } */
+}
+
+void
+h1()
+{
+  omp_allocator_handle_t my_handle;
+  int B1[3]; /* { dg-warning "'my_handle' is used uninitialized" } */
+  /* { dg-warning "variable 'B1' set but not used" "" { target *-*-* } .-1 } */
+  #pragma omp allocate(B1) allocator(my_handle)
+  B1[0] = 5;
+  /* { dg-final { scan-tree-dump-times "__builtin_GOMP_alloc" 1 "gimple" } } */
+  /* { dg-final { scan-tree-dump-times "B1.\[0-9\]+ = __builtin_GOMP_alloc \\(4, 12, my_handle\\);" 1 "gimple" } } */
+  /* { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\(B1.\[0-9\]+, 0B\\);" 1 "gimple" } } */
+}
+
+void
+h2()
+{
+  omp_allocator_handle_t my_handle;
+  int B2[3];  /* { dg-warning "unused variable 'B2'" } */
+  #pragma omp allocate(B2) allocator(my_handle) /* No warning as 'B2' is unused */
+}
+
+void
+h3()
+{
+  omp_allocator_handle_t my_handle;
+  int B3[3] = {1,2,3};  /* { dg-warning "unused variable 'B3'" } */
+  #pragma omp allocate(B3) allocator(my_handle) /* No warning as 'B3' is unused */
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/allocate-11.c b/gcc/testsuite/c-c++-common/gomp/allocate-11.c
index f9ad50abb7f..dceb97f8c5f 100644
--- a/gcc/testsuite/c-c++-common/gomp/allocate-11.c
+++ b/gcc/testsuite/c-c++-common/gomp/allocate-11.c
@@ -10,7 +10,6 @@  f (int i)
   switch (i)  /* { dg-note "switch starts here" } */
     {
       int j;  /* { dg-note "'j' declared here" } */
-      /* { dg-message "sorry, unimplemented: OpenMP 'allocate' directive" "" { target *-*-* } .-1 } */
       #pragma omp allocate(j)
     case 42:  /* { dg-error "switch jumps over OpenMP 'allocate' allocation" } */
       bar ();
@@ -30,9 +29,7 @@  h (int i2)
   return 5;
 
   int k2;  /* { dg-note "'k2' declared here" } */
-  /* { dg-message "sorry, unimplemented: OpenMP 'allocate' directive" "" { target *-*-* } .-1 } */
   int j2 = 4;  /* { dg-note "'j2' declared here" } */
-  /* { dg-message "sorry, unimplemented: OpenMP 'allocate' directive" "" { target *-*-* } .-1 } */
   #pragma omp allocate(k2, j2)
 label:  /* { dg-note "label 'label' defined here" } */
   k2 = 4;
diff --git a/gcc/testsuite/c-c++-common/gomp/allocate-12.c b/gcc/testsuite/c-c++-common/gomp/allocate-12.c
index 3c7c3bb3a2b..1b77db9bd6f 100644
--- a/gcc/testsuite/c-c++-common/gomp/allocate-12.c
+++ b/gcc/testsuite/c-c++-common/gomp/allocate-12.c
@@ -17,7 +17,6 @@  f ()
   omp_allocator_handle_t my_allocator;
   int n = 5;  /* { dg-note "to be allocated variable declared here" } */
   my_allocator = omp_default_mem_alloc; /* { dg-note "modified here" } */
-  /* { dg-message "sorry, unimplemented: OpenMP 'allocate' directive" "" { target *-*-* } .-2 } */
   #pragma omp allocate(n) allocator(my_allocator)  /* { dg-error "variable 'my_allocator' used in the 'allocator' clause must not be modified between declaration of 'n' and its 'allocate' directive" } */
   n = 7;
   return n;
@@ -28,7 +27,6 @@  int
 g ()
 {
   int n = 5;  /* { dg-note "to be allocated variable declared here" } */
-  /* { dg-message "sorry, unimplemented: OpenMP 'allocate' directive" "" { target *-*-* } .-1 } */
   omp_allocator_handle_t my_allocator = omp_low_lat_mem_alloc;  /* { dg-note "declared here" } */
   #pragma omp allocate(n) allocator(my_allocator)  /* { dg-error "variable 'my_allocator' used in the 'allocator' clause must be declared before 'n'" } */
   n = 7;
@@ -42,7 +40,6 @@  h ()
      see gomp/allocate-10.c.  */
   omp_allocator_handle_t my_allocator;
   int n = 5;
-  /* { dg-message "sorry, unimplemented: OpenMP 'allocate' directive" "" { target *-*-* } .-1 } */
   #pragma omp allocate(n) allocator(my_allocator)
   n = 7;
   return n;
diff --git a/gcc/testsuite/c-c++-common/gomp/allocate-15.c b/gcc/testsuite/c-c++-common/gomp/allocate-15.c
index d9600f96c46..15105b9102e 100644
--- a/gcc/testsuite/c-c++-common/gomp/allocate-15.c
+++ b/gcc/testsuite/c-c++-common/gomp/allocate-15.c
@@ -8,7 +8,7 @@  void
 f ()
 {
 
-  int var;  /* { dg-message "sorry, unimplemented: OpenMP 'allocate' directive, used for 'var', not yet supported" } */
+  int var;
   #pragma omp allocate(var)
   var = 5;
 }
@@ -21,7 +21,7 @@  h ()
    #pragma omp parallel
     #pragma omp serial
      {
-       int var2[5];  /* { dg-message "sorry, unimplemented: OpenMP 'allocate' directive, used for 'var2', not yet supported" } */
+       int var2[5];
        #pragma omp allocate(var2)
        var2[0] = 7;
      }
diff --git a/gcc/testsuite/c-c++-common/gomp/allocate-17.c b/gcc/testsuite/c-c++-common/gomp/allocate-17.c
new file mode 100644
index 00000000000..f75af0c2d93
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/allocate-17.c
@@ -0,0 +1,37 @@ 
+/* This file has a syntax error but should not ICE.
+   Namely, a '}' is missing in one(). */
+
+typedef enum omp_allocator_handle_t
+#if __cplusplus >= 201103L
+: __UINTPTR_TYPE__
+#endif
+{
+  omp_default_mem_alloc = 1,
+  omp_low_lat_mem_alloc = 5,
+  __omp_allocator_handle_t_max__ = __UINTPTR_MAX__
+} omp_allocator_handle_t;
+
+#include <stdint.h>
+
+void
+one ()
+{  /* { dg-note "to match this '\{'" "" { target c++ } } */
+  int result = 0, n = 3;
+  #pragma omp target map(tofrom: result) firstprivate(n)
+    {
+      int var = 5; //, var2[n];
+      #pragma omp allocate(var) align(128) allocator(omp_low_lat_mem_alloc) /* { dg-message "sorry, unimplemented: '#pragma omp allocate' not yet supported" "" { target c++ } } */
+       var = 7;
+}
+
+void
+two ()
+{ /* { dg-error "a function-definition is not allowed here before '\{' token" "" { target c++ } } */
+  int scalar = 44;
+  #pragma omp allocate(scalar)
+
+  #pragma omp parallel firstprivate(scalar)
+    scalar = 33;
+}
+/* { dg-error "expected declaration or statement at end of input" "" { target c } .-1 } */
+/* { dg-error "expected '\}' at end of input" "" { target c++ } .-2 } */
diff --git a/gcc/testsuite/c-c++-common/gomp/allocate-9.c b/gcc/testsuite/c-c++-common/gomp/allocate-9.c
index 8e010419a5f..3c11080dd16 100644
--- a/gcc/testsuite/c-c++-common/gomp/allocate-9.c
+++ b/gcc/testsuite/c-c++-common/gomp/allocate-9.c
@@ -86,8 +86,6 @@  int g()
 /* { dg-note "declared here" "" { target c } .-8 } */
 /* { dg-message "sorry, unimplemented: '#pragma omp allocate' not yet supported" "" { target c++ } .-2 } */
     return c2+a2+b2;
-  /* { dg-message "sorry, unimplemented: OpenMP 'allocate' directive" "" { target c } .-5 } */
-  /* { dg-message "sorry, unimplemented: OpenMP 'allocate' directive" "" { target c } .-12 } */
   }
 }
 
diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi
index c6cd825bbaa..8227f89346a 100644
--- a/libgomp/libgomp.texi
+++ b/libgomp/libgomp.texi
@@ -225,7 +225,7 @@  The OpenMP 4.5 specification is fully supported.
 @item Predefined memory spaces, memory allocators, allocator traits
       @tab Y @tab See also @ref{Memory allocation}
 @item Memory management routines @tab Y @tab
-@item @code{allocate} directive @tab N @tab
+@item @code{allocate} directive @tab P @tab Only C, only stack variables
 @item @code{allocate} clause @tab P @tab Initial support
 @item @code{use_device_addr} clause on @code{target data} @tab Y @tab
 @item @code{ancestor} modifier on @code{device} clause @tab Y @tab
@@ -296,7 +296,8 @@  The OpenMP 4.5 specification is fully supported.
 @item Loop transformation constructs @tab N @tab
 @item @code{strict} modifier in the @code{grainsize} and @code{num_tasks}
       clauses of the @code{taskloop} construct @tab Y @tab
-@item @code{align} clause in @code{allocate} directive @tab N @tab
+@item @code{align} clause in @code{allocate} directive @tab P
+      @tab Only C (and only stack variables)
 @item @code{align} modifier in @code{allocate} clause @tab Y @tab
 @item @code{thread_limit} clause to @code{target} construct @tab Y @tab
 @item @code{has_device_addr} clause to @code{target} construct @tab Y @tab
diff --git a/libgomp/testsuite/libgomp.c/allocate-4.c b/libgomp/testsuite/libgomp.c/allocate-4.c
new file mode 100644
index 00000000000..e81cc4093aa
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/allocate-4.c
@@ -0,0 +1,84 @@ 
+/* TODO: move to ../libgomp.c-c++-common once C++ is implemented. */
+/* NOTE: { target c } is unsupported with with the C compiler.  */
+
+/* { dg-do run } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+#include <omp.h>
+#include <stdint.h>
+
+/* { dg-final { scan-tree-dump-times "__builtin_GOMP_alloc \\(" 5 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\(" 5 "gimple" } } */
+
+
+int one ()
+{
+  int sum = 0;
+  #pragma omp allocate(sum)
+  /* { dg-final { scan-tree-dump-times "sum\\.\[0-9\]+ = __builtin_GOMP_alloc \\(4, 4, 0B\\);" 1 "gimple" } } */
+  /* { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\(sum\\.\[0-9\]+, 0B\\);" 1 "gimple" } } */
+
+  /* NOTE: Initializer cannot be omp_init_allocator - as 'A' is
+     in the same scope and the auto-omp_free comes later than
+     any omp_destroy_allocator.  */
+  omp_allocator_handle_t my_allocator = omp_low_lat_mem_alloc;
+  int n = 25;
+  int A[n];
+  #pragma omp allocate(A) align(128) allocator(my_allocator)
+  /* { dg-final { scan-tree-dump-times "A\\.\[0-9\]+ = __builtin_GOMP_alloc \\(128, _\[0-9\]+, my_allocator\\);" 1 "gimple" } } */
+  /* { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\(A\\.\[0-9\]+, 0B\\);" 1 "gimple" } } */
+
+  if (((intptr_t)A) % 128 != 0)
+    __builtin_abort ();
+  for (int i = 0; i < n; ++i)
+    A[i] = i;
+
+  omp_alloctrait_t traits[1] = { { omp_atk_alignment, 64 } };
+  my_allocator = omp_init_allocator(omp_low_lat_mem_space,1,traits);
+  {
+    int B[n] = { };
+    int C[5] = {1,2,3,4,5};
+    #pragma omp allocate(B,C) allocator(my_allocator)
+    /* { dg-final { scan-tree-dump-times "B\\.\[0-9\]+ = __builtin_GOMP_alloc \\(\[0-9\]+, _\[0-9\]+, my_allocator\\);" 1 "gimple" } } */
+    /* { dg-final { scan-tree-dump-times "C\\.\[0-9\]+ = __builtin_GOMP_alloc \\(\[0-9\]+, 20, my_allocator\\);" 1 "gimple" } } */
+    /* { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\(B\\.\[0-9\]+, 0B\\);" 1 "gimple" } } */
+    /* { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\(C\\.\[0-9\]+, 0B\\);" 1 "gimple" } } */
+
+    int D[5] = {11,22,33,44,55};
+    #pragma omp allocate(D) align(256)
+    /* { dg-final { scan-tree-dump-times "D\\.\[0-9\]+ = __builtin_GOMP_alloc \\(256, 20, 0B\\);" 1 "gimple" } } */
+    /* { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\(D\\.\[0-9\]+, 0B\\);" 1 "gimple" } } */
+
+    if (((intptr_t) B) % 64 != 0)
+      __builtin_abort ();
+    if (((intptr_t) C) % 64 != 0)
+      __builtin_abort ();
+    if (((intptr_t) D) % 64 != 0)
+      __builtin_abort ();
+
+    for (int i = 0; i < 5; ++i)
+      {
+	if (C[i] != i+1)
+	  __builtin_abort ();
+	if (D[i] != i+1 + 10*(i+1))
+	  __builtin_abort ();
+      }
+
+    for (int i = 0; i < n; ++i)
+      {
+	if (B[i] != 0)
+	  __builtin_abort ();
+	sum += A[i]+B[i]+C[i%5]+D[i%5];
+      }
+  }
+  omp_destroy_allocator (my_allocator);
+  return sum;
+}
+
+int
+main ()
+{
+  if (one () != 1200)
+    __builtin_abort ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/allocate-5.c b/libgomp/testsuite/libgomp.c/allocate-5.c
new file mode 100644
index 00000000000..beaf16440e1
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/allocate-5.c
@@ -0,0 +1,126 @@ 
+/* TODO: move to ../libgomp.c-c++-common once C++ is implemented. */
+/* NOTE: { target c } is unsupported with with the C compiler.  */
+
+/* { dg-do run } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+#include <omp.h>
+#include <stdint.h>
+
+/* { dg-final { scan-tree-dump-not "__builtin_stack_save" "gimple" } } */
+/* { dg-final { scan-tree-dump-not "__builtin_alloca" "gimple" } } */
+/* { dg-final { scan-tree-dump-not "__builtin_stack_restore" "gimple" } } */
+
+/* { dg-final { scan-tree-dump-times "__builtin_GOMP_alloc \\(" 5 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\(" 5 "gimple" } } */
+
+void
+one ()
+{
+  int result = 0, n = 3;
+  #pragma omp target map(tofrom: result) firstprivate(n)
+    {
+      int var = 5, var2[n];
+      #pragma omp allocate(var,var2) align(128) allocator(omp_low_lat_mem_alloc)
+/* { dg-final { scan-tree-dump-times "var\\.\[0-9\]+ = __builtin_GOMP_alloc \\(128, 4, 5\\);" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "var2\\.\[0-9\]+ = __builtin_GOMP_alloc \\(128, D\\.\[0-9\]+, 5\\);" 1 "gimple" } } */
+
+/* { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\(var\\.\[0-9\]+, 0B\\);" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\(var2\\.\[0-9\]+, 0B\\);" 1 "gimple" } } */
+
+      if ((intptr_t) &var % 128 != 0)
+	__builtin_abort ();
+      if ((intptr_t) var2 % 128 != 0)
+	__builtin_abort ();
+      if (var != 5)
+	__builtin_abort ();
+
+      #pragma omp parallel for
+      for (int i = 0; i < n; ++i)
+	var2[i] = (i+33);
+
+      #pragma omp loop reduction(+:result)
+      for (int i = 0; i < n; ++i)
+	result += var + var2[i];
+    }
+  if (result != (3*5 + 33 + 34 + 35))
+    __builtin_abort ();
+}
+
+void
+two ()
+{
+  struct st {
+    int a, b;
+  };
+  int scalar = 44, array[5] = {1,2,3,4,5};
+  struct st s = {.a=11, .b=56};
+  #pragma omp allocate(scalar, array, s)
+/* { dg-final { scan-tree-dump-times "scalar\\.\[0-9\]+ = __builtin_GOMP_alloc \\(4, 4, 0B\\);" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "array\\.\[0-9\]+ = __builtin_GOMP_alloc \\(4, 20, 0B\\);" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "s\\.\[0-9\]+ = __builtin_GOMP_alloc \\(4, 8, 0B\\);" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\(scalar\\.\[0-9\]+, 0B\\);" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\(array\\.\[0-9\]+, 0B\\);" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\(s\\.\[0-9\]+, 0B\\);" 1 "gimple" } } */
+
+  #pragma omp parallel firstprivate(scalar) firstprivate(array) firstprivate(s)
+  {
+    if (scalar != 44)
+      __builtin_abort ();
+    scalar = 33;
+    for (int i = 0; i < 5; ++i)
+      if (array[i] != i+1)
+	__builtin_abort ();
+    for (int i = 0; i < 5; ++i)
+      array[i] = 10*(i+1);
+    if (s.a != 11 || s.b != 56)
+      __builtin_abort ();
+    s.a = 74;
+    s.b = 674;
+  }
+  if (scalar != 44)
+    __builtin_abort ();
+  for (int i = 0; i < 5; ++i)
+    if (array[i] != i+1)
+      __builtin_abort ();
+  if (s.a != 11 || s.b != 56)
+    __builtin_abort ();
+
+  #pragma omp target defaultmap(firstprivate : scalar) defaultmap(none : aggregate) defaultmap(none : pointer)
+  {
+    if (scalar != 44)
+      __builtin_abort ();
+    scalar = 33;
+  }
+  if (scalar != 44)
+    __builtin_abort ();
+
+  #pragma omp target defaultmap(none : scalar) defaultmap(firstprivate : aggregate) defaultmap(none : pointer)
+  {
+    for (int i = 0; i < 5; ++i)
+      if (array[i] != i+1)
+	__builtin_abort ();
+    for (int i = 0; i < 5; ++i)
+      array[i] = 10*(i+1);
+  }
+  for (int i = 0; i < 5; ++i)
+    if (array[i] != i+1)
+      __builtin_abort ();
+  #pragma omp target defaultmap(none : scalar) defaultmap(firstprivate : aggregate) defaultmap(none : pointer)
+  {
+    if (s.a != 11 || s.b != 56)
+      __builtin_abort ();
+    s.a = 74;
+    s.b = 674;
+  }
+  if (s.a != 11 || s.b != 56)
+    __builtin_abort ();
+}
+
+int
+main ()
+{
+  one ();
+  two ();
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/allocate-6.c b/libgomp/testsuite/libgomp.c/allocate-6.c
new file mode 100644
index 00000000000..6d7278ce571
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/allocate-6.c
@@ -0,0 +1,319 @@ 
+/* TODO: move to ../libgomp.c-c++-common once C++ is implemented. */
+/* NOTE: { target c } is unsupported with with the C compiler.  */
+
+/* { dg-do run } */
+/* { dg-additional-options "-fdump-tree-omplower" } */
+
+/* For the 4 vars in omp_parallel, 4 in omp_target and 1 of 2 in no_alloc2_func.  */
+/* { dg-final { scan-tree-dump-times "__builtin_GOMP_alloc \\(" 9 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "__builtin_GOMP_free \\(" 9 "omplower" } } */
+
+#include <omp.h>
+
+void
+check_int (int *x, int y)
+{
+  if (*x != y)
+    __builtin_abort ();
+}
+
+void
+check_ptr (int **x, int *y)
+{
+  if (*x != y)
+    __builtin_abort ();
+}
+
+
+int
+no_alloc_func ()
+{
+  /* There is no __builtin_GOMP_alloc / __builtin_GOMP_free as
+     allocator == omp_default_mem_alloc (known at compile time. */
+  int no_alloc;
+  #pragma omp allocate(no_alloc) allocator(omp_default_mem_alloc)
+  no_alloc = 7;
+  return no_alloc;
+}
+
+int
+no_alloc2_func()
+{
+  /* There is no __builtin_GOMP_alloc / __builtin_GOMP_free as
+     no_alloc2 is TREE_UNUSED.  But there is for is_alloc2.  */
+  int no_alloc2, is_alloc2;
+  #pragma omp allocate(no_alloc2, is_alloc2)
+  is_alloc2 = 7;
+  return is_alloc2;
+}
+
+
+void
+omp_parallel ()
+{
+  int n = 6;
+  int iii = 5, jjj[5], kkk[n];
+  int *ptr = (int *) 0x1234;
+  #pragma omp allocate(iii, jjj, kkk, ptr)
+
+  for (int i = 0; i < 5; i++)
+    jjj[i] = 3*i;
+  for (int i = 0; i < 6; i++)
+    kkk[i] = 7*i;
+
+  #pragma omp parallel default(none) firstprivate(iii, jjj, kkk, ptr) if(0)
+  {
+    if (iii != 5)
+      __builtin_abort();
+    iii = 7;
+    check_int (&iii, 7);
+    for (int i = 0; i < 5; i++)
+      if (jjj[i] != 3*i)
+	__builtin_abort ();
+    for (int i = 0; i < 6; i++)
+      if (kkk[i] != 7*i)
+	__builtin_abort ();
+    for (int i = 0; i < 5; i++)
+      jjj[i] = 4*i;
+    for (int i = 0; i < 6; i++)
+      kkk[i] = 8*i;
+    for (int i = 0; i < 5; i++)
+      check_int (&jjj[i], 4*i);
+    for (int i = 0; i < 6; i++)
+      check_int (&kkk[i], 8*i);
+    if (ptr != (int *) 0x1234)
+      __builtin_abort ();
+    ptr = (int *) 0xabcd;
+    if (ptr != (int *) 0xabcd)
+      __builtin_abort ();
+    check_ptr (&ptr, (int *) 0xabcd);
+  }
+  if (iii != 5)
+    __builtin_abort ();
+  check_int (&iii, 5);
+  for (int i = 0; i < 5; i++)
+    {
+      if (jjj[i] != 3*i)
+	__builtin_abort ();
+      check_int (&jjj[i], 3*i);
+    }
+  for (int i = 0; i < 6; i++)
+    {
+      if (kkk[i] != 7*i)
+	__builtin_abort ();
+      check_int (&kkk[i], 7*i);
+    }
+  if (ptr != (int *) 0x1234)
+    __builtin_abort ();
+  check_ptr (&ptr, (int *) 0x1234);
+
+  #pragma omp parallel default(firstprivate) if(0)
+  {
+    if (iii != 5)
+      __builtin_abort();
+    iii = 7;
+    check_int (&iii, 7);
+    for (int i = 0; i < 5; i++)
+      if (jjj[i] != 3*i)
+	__builtin_abort ();
+    for (int i = 0; i < 6; i++)
+      if (kkk[i] != 7*i)
+	__builtin_abort ();
+    for (int i = 0; i < 5; i++)
+      jjj[i] = 4*i;
+    for (int i = 0; i < 6; i++)
+      kkk[i] = 8*i;
+    for (int i = 0; i < 5; i++)
+      check_int (&jjj[i], 4*i);
+    for (int i = 0; i < 6; i++)
+      check_int (&kkk[i], 8*i);
+    if (ptr != (int *) 0x1234)
+      __builtin_abort ();
+    ptr = (int *) 0xabcd;
+    if (ptr != (int *) 0xabcd)
+      __builtin_abort ();
+    check_ptr (&ptr, (int *) 0xabcd);
+  }
+  if (iii != 5)
+    __builtin_abort ();
+  check_int (&iii, 5);
+  for (int i = 0; i < 5; i++)
+    {
+      if (jjj[i] != 3*i)
+	__builtin_abort ();
+      check_int (&jjj[i], 3*i);
+    }
+  for (int i = 0; i < 6; i++)
+    {
+      if (kkk[i] != 7*i)
+	__builtin_abort ();
+      check_int (&kkk[i], 7*i);
+    }
+  if (ptr != (int *) 0x1234)
+    __builtin_abort ();
+  check_ptr (&ptr, (int *) 0x1234);
+}
+
+
+
+void
+omp_target ()
+{
+  int n = 6;
+  int iii = 5, jjj[5], kkk[n];
+  int *ptr = (int *) 0x1234;
+  #pragma omp allocate(iii, jjj, kkk, ptr)
+
+  for (int i = 0; i < 5; i++)
+    jjj[i] = 3*i;
+  for (int i = 0; i < 6; i++)
+    kkk[i] = 7*i;
+
+  #pragma omp target defaultmap(none) firstprivate(iii, jjj, kkk, ptr)
+  {
+    if (iii != 5)
+      __builtin_abort();
+    iii = 7;
+    check_int (&iii, 7);
+    for (int i = 0; i < 5; i++)
+      if (jjj[i] != 3*i)
+	__builtin_abort ();
+    for (int i = 0; i < 6; i++)
+      if (kkk[i] != 7*i)
+	__builtin_abort ();
+    for (int i = 0; i < 5; i++)
+      jjj[i] = 4*i;
+    for (int i = 0; i < 6; i++)
+      kkk[i] = 8*i;
+    for (int i = 0; i < 5; i++)
+      check_int (&jjj[i], 4*i);
+    for (int i = 0; i < 6; i++)
+      check_int (&kkk[i], 8*i);
+    if (ptr != (int *) 0x1234)
+      __builtin_abort ();
+    ptr = (int *) 0xabcd;
+    if (ptr != (int *) 0xabcd)
+      __builtin_abort ();
+    check_ptr (&ptr, (int *) 0xabcd);
+  }
+  if (iii != 5)
+    __builtin_abort ();
+  check_int (&iii, 5);
+  for (int i = 0; i < 5; i++)
+    {
+      if (jjj[i] != 3*i)
+	__builtin_abort ();
+      check_int (&jjj[i], 3*i);
+    }
+  for (int i = 0; i < 6; i++)
+    {
+      if (kkk[i] != 7*i)
+	__builtin_abort ();
+      check_int (&kkk[i], 7*i);
+    }
+  if (ptr != (int *) 0x1234)
+    __builtin_abort ();
+  check_ptr (&ptr, (int *) 0x1234);
+
+  #pragma omp target defaultmap(firstprivate)
+  {
+    if (iii != 5)
+      __builtin_abort();
+    iii = 7;
+    check_int (&iii, 7);
+    for (int i = 0; i < 5; i++)
+      if (jjj[i] != 3*i)
+	__builtin_abort ();
+    for (int i = 0; i < 6; i++)
+      if (kkk[i] != 7*i)
+	__builtin_abort ();
+    for (int i = 0; i < 5; i++)
+      jjj[i] = 4*i;
+    for (int i = 0; i < 6; i++)
+      kkk[i] = 8*i;
+    for (int i = 0; i < 5; i++)
+      check_int (&jjj[i], 4*i);
+    for (int i = 0; i < 6; i++)
+      check_int (&kkk[i], 8*i);
+    if (ptr != (int *) 0x1234)
+      __builtin_abort ();
+    ptr = (int *) 0xabcd;
+    if (ptr != (int *) 0xabcd)
+      __builtin_abort ();
+    check_ptr (&ptr, (int *) 0xabcd);
+  }
+  if (iii != 5)
+    __builtin_abort ();
+  check_int (&iii, 5);
+  for (int i = 0; i < 5; i++)
+    {
+      if (jjj[i] != 3*i)
+	__builtin_abort ();
+      check_int (&jjj[i], 3*i);
+    }
+  for (int i = 0; i < 6; i++)
+    {
+      if (kkk[i] != 7*i)
+	__builtin_abort ();
+      check_int (&kkk[i], 7*i);
+    }
+  if (ptr != (int *) 0x1234)
+    __builtin_abort ();
+  check_ptr (&ptr, (int *) 0x1234);
+
+  #pragma omp target defaultmap(tofrom)
+  {
+    if (iii != 5)
+      __builtin_abort();
+    iii = 7;
+    check_int (&iii, 7);
+    for (int i = 0; i < 5; i++)
+      if (jjj[i] != 3*i)
+	__builtin_abort ();
+    for (int i = 0; i < 6; i++)
+      if (kkk[i] != 7*i)
+	__builtin_abort ();
+    for (int i = 0; i < 5; i++)
+      jjj[i] = 4*i;
+    for (int i = 0; i < 6; i++)
+      kkk[i] = 8*i;
+    for (int i = 0; i < 5; i++)
+      check_int (&jjj[i], 4*i);
+    for (int i = 0; i < 6; i++)
+      check_int (&kkk[i], 8*i);
+    if (ptr != (int *) 0x1234)
+      __builtin_abort ();
+    ptr = (int *) 0xabcd;
+    if (ptr != (int *) 0xabcd)
+      __builtin_abort ();
+    check_ptr (&ptr, (int *) 0xabcd);
+  }
+
+  if (iii != 7)
+    __builtin_abort ();
+  check_int (&iii, 7);
+  for (int i = 0; i < 5; i++)
+    {
+      if (jjj[i] != 4*i)
+	__builtin_abort ();
+      check_int (&jjj[i], 4*i);
+    }
+  for (int i = 0; i < 6; i++)
+    {
+      if (kkk[i] != 8*i)
+	__builtin_abort ();
+      check_int (&kkk[i], 8*i);
+    }
+  if (ptr != (int *) 0xabcd)
+    __builtin_abort ();
+  check_ptr (&ptr, (int *) 0xabcd);
+}
+
+
+int
+main ()
+{
+  omp_parallel ();
+  omp_target ();
+  return 0;
+}