Re: [gomp4.1] Various accelerator updates from OpenMP 4.1

[Jakub Jelinek <jakub@redhat.com>]

On Fri, Sep 04, 2015 at 09:07:02PM +0300, Ilya Verbin wrote:
> It seems that there is a bug some here:
> 
> Here is the reproducer:
> 
> #pragma omp declare target
> int a[1];
> #pragma omp end declare target
> 
> void foo ()
> {
>   #pragma omp target map(to: a[0:1])
>     a;
> }
> 
> 
> lookup_decl (var, ctx) tries to lookup for 'a', but ctx->cb.decl_map->get ()
> returns null-pointer.

Fixed thusly, tested on x86_64-linux, committed to gomp4.1 branch.

2015-09-07  Jakub Jelinek  <jakub@redhat.com>

	* omp-low.c (scan_sharing_clauses): Don't ignore map with
	declare target vars for GOMP_MAP_FIRSTPRIVATE_POINTER,
	unless the decl is an array.
	(lower_omp_target): Ignore GOMP_MAP_FIRSTPRIVATE_POINTER map with
	declare target var if it is an array.

	* testsuite/libgomp.c/target-26.c: New test.

--- gcc/omp-low.c.jj	2015-09-04 11:34:45.000000000 +0200
+++ gcc/omp-low.c	2015-09-07 14:10:36.198517500 +0200
@@ -2060,6 +2060,8 @@ scan_sharing_clauses (tree clauses, omp_
 	     directly.  */
 	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 	      && DECL_P (decl)
+	      && (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER
+		  || TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
 	      && is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
 	      && varpool_node::get_create (decl)->offloadable)
 	    break;
@@ -2284,6 +2286,8 @@ scan_sharing_clauses (tree clauses, omp_
 	    break;
 	  decl = OMP_CLAUSE_DECL (c);
 	  if (DECL_P (decl)
+	      && (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER
+		  || TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
 	      && is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
 	      && varpool_node::get_create (decl)->offloadable)
 	    break;
@@ -13358,6 +13362,10 @@ lower_omp_target (gimple_stmt_iterator *
 	  {
 	    if (TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
 	      {
+		if (is_global_var (maybe_lookup_decl_in_outer_ctx (var, ctx))
+		    && varpool_node::get_create (var)->offloadable)
+		  continue;
+
 		tree type = build_pointer_type (TREE_TYPE (var));
 		tree new_var = lookup_decl (var, ctx);
 		x = create_tmp_var_raw (type, get_name (new_var));
@@ -14081,6 +14089,12 @@ lower_omp_target (gimple_stmt_iterator *
 		HOST_WIDE_INT offset = 0;
 		gcc_assert (prev);
 		var = OMP_CLAUSE_DECL (c);
+		if (DECL_P (var)
+		    && TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE
+		    && is_global_var (maybe_lookup_decl_in_outer_ctx (var,
+								      ctx))
+		    && varpool_node::get_create (var)->offloadable)
+		  break;
 		if (TREE_CODE (var) == INDIRECT_REF
 		    && TREE_CODE (TREE_OPERAND (var, 0)) == COMPONENT_REF)
 		  var = TREE_OPERAND (var, 0);
--- libgomp/testsuite/libgomp.c/target-26.c.jj	2015-09-07 11:59:08.665425993 +0200
+++ libgomp/testsuite/libgomp.c/target-26.c	2015-09-07 12:38:56.000000000 +0200
@@ -0,0 +1,36 @@
+extern void abort (void);
+#pragma omp declare target
+int a[4] = { 2, 3, 4, 5 }, *b;
+#pragma omp end declare target
+
+int
+main ()
+{
+  int err;
+  int c[3] = { 6, 7, 8 };
+  b = c;
+  #pragma omp target map(to: a[0:2], b[0:2]) map(from: err)
+  err = a[0] != 2 || a[1] != 3 || a[2] != 4 || a[3] != 5 || b[0] != 6 || b[1] != 7;
+  if (err)
+    abort ();
+  a[1] = 9;
+  a[2] = 10;
+  #pragma omp target map(always,to:a[1:2]) map(from: err)
+  err = a[0] != 2 || a[1] != 9 || a[2] != 10 || a[3] != 5;
+  if (err)
+    abort ();
+  #pragma omp parallel firstprivate(a, b, c, err) num_threads (2)
+  #pragma omp single
+  {
+    b = c + 1;
+    a[0] = 11;
+    a[2] = 13;
+    c[1] = 14;
+    int d = 0;
+    #pragma omp target map(to: a[0:3], b[d:2]) map (from: err)
+    err = a[0] != 11 || a[1] != 9 || a[2] != 13 || b[0] != 14 || b[1] != 8;
+    if (err)
+      abort ();
+  }
+  return 0;
+}

	Jakub