public inbox for gcc-patches@gcc.gnu.org
 help / color / mirror / Atom feed
* [Patch][Fortran/OpenMP] Don't create "alloc:" for 'target exit data'
@ 2019-10-18 10:13 Tobias Burnus
  2019-10-24  7:29 ` *ping* " Tobias Burnus
  2019-10-30 10:37 ` Jakub Jelinek
  0 siblings, 2 replies; 7+ messages in thread
From: Tobias Burnus @ 2019-10-18 10:13 UTC (permalink / raw)
  To: gcc-patches, fortran, Jakub Jelinek

[-- Attachment #1: Type: text/plain, Size: 1081 bytes --]

Currently, one has for
   !$omp target exit data map(delete:x)
in the original dump:
   #pragma omp target exit data map(delete:*x) map(alloc:x [pointer 
assign, bias: 0])

The "alloc:" not only does not make sense but also gives run-time 
messages like:
libgomp: GOMP_target_enter_exit_data unhandled kind 0x04

[Depending on the data type, in gfc_trans_omp_clauses's OMP_LIST_MAP, 
add map clauses of type GOMP_MAP_POINTER and/or GOMP_MAP_TO_PSET.]

That's for release:/delete:. However, for 'target exit data' 
(GOMP_target_enter_exit_data) the same issue occurs for "from:"/"always, 
from:".  But "from:" implies "alloc:". – While "alloc:" does not make 
sense for "target exit data" or "update", for "target" or "target data" 
it surely matters. Hence, I only exclude "from:" for exit data and update.

See attached patch. I have additionally Fortran-fied 
libgomp.c/target-20.c to have at least one 'enter/exit target data' test 
case for Fortran.

Build + regtested on x86_64-gnu-linux w/o offloading. And I have tested 
the new test case with nvptx.

Tobias


[-- Attachment #2: exit-target-data-v2.diff --]
[-- Type: text/x-patch, Size: 7945 bytes --]

 	gcc/fortran/
	* trans-openmp.c (gfc_trans_omp_clauses): Do not create
	map(alloc:) for map(delete:/release:) and for
	(from:/always,from:) only if new arg require_from_alloc is true,
	which is the default.
	(gfc_trans_omp_target_exit_data, gfc_trans_omp_target_update):
	Call it with require_from_alloc = false.

	libgomp/
	* testsuite/libgomp.fortran/target9.f90: New.

diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index dad11a24430..f890629c73d 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -1852,7 +1852,8 @@ static vec<tree, va_heap, vl_embed> *doacross_steps;
 
 static tree
 gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
-		       locus where, bool declare_simd = false)
+		       locus where, bool declare_simd = false,
+		       bool require_from_alloc = true)
 {
   tree omp_clauses = NULL_TREE, chunk_size, c;
   int list, ifc;
@@ -2163,6 +2164,16 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 	      if (!n->sym->attr.referenced)
 		continue;
 
+	      /* map(alloc:) etc. is not needed for delete/release
+		 For 'from:', it is needed when setting up the environment
+		 but not for updating or copying out of the data.  */
+	      bool no_extra_pointer = n->u.map_op == OMP_MAP_DELETE
+				      || n->u.map_op == OMP_MAP_RELEASE
+				      || (!require_from_alloc
+					  && (n->u.map_op == OMP_MAP_FROM
+					      || n->u.map_op
+						     == OMP_MAP_ALWAYS_FROM));
+
 	      tree node = build_omp_clause (input_location, OMP_CLAUSE_MAP);
 	      tree node2 = NULL_TREE;
 	      tree node3 = NULL_TREE;
@@ -2172,7 +2183,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		TREE_ADDRESSABLE (decl) = 1;
 	      if (n->expr == NULL || n->expr->ref->u.ar.type == AR_FULL)
 		{
-		  if (POINTER_TYPE_P (TREE_TYPE (decl))
+		  if (!no_extra_pointer
+		      && POINTER_TYPE_P (TREE_TYPE (decl))
 		      && (gfc_omp_privatize_by_reference (decl)
 			  || GFC_DECL_GET_SCALAR_POINTER (decl)
 			  || GFC_DECL_GET_SCALAR_ALLOCATABLE (decl)
@@ -2208,17 +2220,20 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 					  ptr);
 		      ptr = build_fold_indirect_ref (ptr);
 		      OMP_CLAUSE_DECL (node) = ptr;
-		      node2 = build_omp_clause (input_location,
-						OMP_CLAUSE_MAP);
-		      OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET);
-		      OMP_CLAUSE_DECL (node2) = decl;
-		      OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
-		      node3 = build_omp_clause (input_location,
-						OMP_CLAUSE_MAP);
-		      OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER);
-		      OMP_CLAUSE_DECL (node3)
-			= gfc_conv_descriptor_data_get (decl);
-		      OMP_CLAUSE_SIZE (node3) = size_int (0);
+		      if (!no_extra_pointer)
+			{
+			  node2 = build_omp_clause (input_location,
+						    OMP_CLAUSE_MAP);
+			  OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET);
+			  OMP_CLAUSE_DECL (node2) = decl;
+			  OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
+			  node3 = build_omp_clause (input_location,
+						    OMP_CLAUSE_MAP);
+			  OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER);
+			  OMP_CLAUSE_DECL (node3)
+				= gfc_conv_descriptor_data_get (decl);
+			  OMP_CLAUSE_SIZE (node3) = size_int (0);
+			}
 
 		      /* We have to check for n->sym->attr.dimension because
 			 of scalar coarrays.  */
@@ -2302,6 +2317,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 				      ptr);
 		  OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (ptr);
 
+		  if (no_extra_pointer)
+		    goto skip_extra_map_pointer;
+
 		  if (POINTER_TYPE_P (TREE_TYPE (decl))
 		      && GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (TREE_TYPE (decl))))
 		    {
@@ -2346,6 +2364,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		  OMP_CLAUSE_SIZE (node3)
 		    = fold_build2 (MINUS_EXPR, sizetype, ptr, ptr2);
 		}
+
+	      skip_extra_map_pointer:
+
 	      switch (n->u.map_op)
 		{
 		case OMP_MAP_ALLOC:
@@ -4979,7 +5000,7 @@ gfc_trans_omp_target_exit_data (gfc_code *code)
 
   gfc_start_block (&block);
   omp_clauses = gfc_trans_omp_clauses (&block, code->ext.omp_clauses,
-				       code->loc);
+				       code->loc, false, false);
   stmt = build1_loc (input_location, OMP_TARGET_EXIT_DATA, void_type_node,
 		     omp_clauses);
   gfc_add_expr_to_block (&block, stmt);
@@ -4994,7 +5015,7 @@ gfc_trans_omp_target_update (gfc_code *code)
 
   gfc_start_block (&block);
   omp_clauses = gfc_trans_omp_clauses (&block, code->ext.omp_clauses,
-				       code->loc);
+				       code->loc, false, false);
   stmt = build1_loc (input_location, OMP_TARGET_UPDATE, void_type_node,
 		     omp_clauses);
   gfc_add_expr_to_block (&block, stmt);
diff --git a/libgomp/testsuite/libgomp.fortran/target9.f90 b/libgomp/testsuite/libgomp.fortran/target9.f90
new file mode 100644
index 00000000000..91d60a33307
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target9.f90
@@ -0,0 +1,123 @@
+! { dg-require-effective-target offload_device_nonshared_as } */
+
+module target_test
+  implicit none (type, external)
+  integer, parameter :: N = 40
+  integer :: sum
+  integer :: var1 = 1
+  integer :: var2 = 2
+
+  !$omp declare target to(D)
+  integer :: D(N) = 0
+contains
+  subroutine enter_data (X)
+    integer :: X(:)
+    !$omp target enter data map(to: var1, var2, X) map(alloc: sum)
+  end subroutine enter_data
+
+  subroutine exit_data_0 (D)
+    integer :: D(N)
+    !$omp target exit data map(delete: D)
+  end subroutine exit_data_0
+
+  subroutine exit_data_1 ()
+    !$omp target exit data map(from: var1)
+  end subroutine exit_data_1
+
+  subroutine exit_data_2 (X)
+    integer :: X(N)
+    !$omp target exit data map(from: var2) map(release: X, sum)
+  end subroutine exit_data_2
+
+  subroutine exit_data_3 (p, idx)
+    integer :: p(:)
+    integer, value :: idx
+    !$omp target exit data map(from: p(idx))
+  end subroutine exit_data_3
+
+  subroutine test_nested ()
+    integer :: X, Y, Z
+    X = 0
+    Y = 0
+    Z = 0
+
+    !$omp target data map(from: X, Y, Z)
+      !$omp target data map(from: X, Y, Z)
+        !$omp target map(from: X, Y, Z)
+          X = 1337
+          Y = 1337
+          Z = 1337
+        !$omp end target
+        if (X /= 0) stop 11
+        if (Y /= 0) stop 12
+        if (Z /= 0) stop 13
+
+        !$omp target exit data map(from: X) map(release: Y)
+        if (X /= 0) stop 14
+        if (Y /= 0) stop 15
+
+        !$omp target exit data map(release: Y) map(delete: Z)
+        if (Y /= 0) stop 16
+        if (Z /= 0) stop 17
+      !$omp end target data
+      if (X /= 1337) stop 18
+      if (Y /= 0) stop 19
+      if (Z /= 0) stop 20
+
+      !$omp target map(from: X)
+        X = 2448
+      !$omp end target
+      if (X /= 2448) stop 21
+      if (Y /= 0) stop 22
+      if (Z /= 0) stop 23
+
+      X = 4896
+    !$omp end target data
+    if (X /= 4896) stop 24
+    if (Y /= 0) stop 25
+    if (Z /= 0) stop 26
+  end subroutine test_nested
+end module target_test
+
+program main
+  use target_test
+  implicit none (type, external)
+
+  integer, allocatable :: X(:)
+  integer, pointer, contiguous :: Y(:)
+
+
+  allocate(X(N), Y(N))
+  X(10) = 10
+  Y(20) = 20
+  call enter_data (X)
+
+  call exit_data_0 (D)  ! This should have no effect on D.
+
+  !$omp target map(alloc: var1, var2, X) map(to: Y) map(always, from: sum)
+    var1 = var1 + X(10)
+    var2 = var2 + Y(20)
+    sum = var1 + var2
+    D(sum) = D(sum) + 1
+  !$omp end target
+
+  if (var1 /= 1) stop 1
+  if (var2 /= 2) stop 2
+  if (sum /= 33) stop 3
+
+  call exit_data_1 ()
+  if (var1 /= 11) stop 4
+  if (var2 /= 2) stop 5
+
+  ! Increase refcount of already mapped X(1:N).
+  !$omp target enter data map(alloc: X(16:17))
+
+  call exit_data_2 (X)
+  if (var2 /= 22) stop 6
+
+  call exit_data_3 (X, 5) ! Unmap X(1:N).
+
+  deallocate (X, Y)
+
+  call test_nested ()
+end program main

^ permalink raw reply	[flat|nested] 7+ messages in thread

end of thread, other threads:[~2019-11-11  9:16 UTC | newest]

Thread overview: 7+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2019-10-18 10:13 [Patch][Fortran/OpenMP] Don't create "alloc:" for 'target exit data' Tobias Burnus
2019-10-24  7:29 ` *ping* " Tobias Burnus
2019-10-30 10:37 ` Jakub Jelinek
2019-10-30 15:50   ` Tobias Burnus
2019-10-30 15:58     ` Jakub Jelinek
2019-10-30 16:46       ` Tobias Burnus
2019-11-11  9:39     ` Thomas Schwinge

This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).