[1/4,OpenACC] Attach / Detach generic infrastructure

Message ID 74507fb6-a3ef-2e24-c24a-fe18ba726e5f@codesourcery.com
State New
Headers show
Series
  • [1/4,OpenACC] Attach / Detach generic infrastructure
Related show

Commit Message

Cesar Philippidis Oct. 26, 2018, 4:39 a.m.
This patch series adds support for the new attach / detach clauses
introduced in OpenACC 2.6 to the C and C++ front ends. Julian is
working patches for the Fortran front end along with the runtime.

As their names somewhat imply, attach and detach are new data clauses
that are used to support manual deep copy in OpenACC. Specifically,
OpenACC 2.6 allows users to specify individual structure fields inside
data clauses, whereas before that would only work inside the update
directive. The attach and detach clauses allow users to update the
pointers in structure fields with their on-device counterparts.

As an example, consider the the following code:

  struct { int *a, b } s;
  int *z = ...

  #pragma acc enter data copyin(a[:N], s)
  ...
  s.a = z;
  #pragma acc enter data attach(s.a)
  ...
  #pragma acc exit data detach(s.a)
  #pragma acc exit data copyout(s)

Because the attach clause updates field s.a with the device address,
"acc exit data detach" must be used to restore the host pointer
contents before that value is copied back to the host.

This patch in particular adds the generic infrastructure for the attach
and detach clauses. All of the front ends lower the attach clause as
GOMP_MAP_DETACH data mapping. However, if a detachment is finalized, e.g.

  #pragma acc exit data finalize detach(ptr)

the gimplifier will promote it to GOMP_MAP_FORCE_FINALIZE. Also, this
patch teaches the gimplifier how to ignore GOMP_MAP_STRUCT for the
target update constructs.

Is this patch OK for trunk? I bootstrapped and regression tested it
for x86_64 Linux with nvptx offloading.

Thanks,
Cesar

Patch

2018-XX-YY  Cesar Philippidis  <cesar@codesourcery.com>

	gcc/
	* gimplify.c (gimplify_adjust_omp_clauses): Filter out
	GOMP_MAP_STRUCT for acc exit data.
	(gimplify_omp_target_update): Promote GOMP_MAP_DETACH
	to GOMP_MAP_FORCE_DETACH when the finalize clause is present.
	* omp-low.c (lower_omp_target): Add support for GOMP_MAP_{ATTACH,
	DETACH, FORCE_DETACH}.
	* tree-pretty-print.c (dump_omp_clause): Likewise.

	gcc/c-family/
	* c-pragma.h (enum pragma_omp_clause): Define
	PRAGMA_OACC_CLAUSE_{ATTACH,DETACH}.

	include/
	* gomp-constants.h (GOMP_MAP_DEEP_COPY): Define.
	(enum gomp_map_kind): Add GOMP_MAP_{ATTACH, DETACH, FORCE_DETACH}.


diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index b322547b11a..ab4c03b21f1 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -131,11 +131,13 @@  enum pragma_omp_clause {
 
   /* Clauses for OpenACC.  */
   PRAGMA_OACC_CLAUSE_ASYNC,
+  PRAGMA_OACC_CLAUSE_ATTACH,
   PRAGMA_OACC_CLAUSE_AUTO,
   PRAGMA_OACC_CLAUSE_COPY,
   PRAGMA_OACC_CLAUSE_COPYOUT,
   PRAGMA_OACC_CLAUSE_CREATE,
   PRAGMA_OACC_CLAUSE_DELETE,
+  PRAGMA_OACC_CLAUSE_DETACH,
   PRAGMA_OACC_CLAUSE_DEVICEPTR,
   PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT,
   PRAGMA_OACC_CLAUSE_FINALIZE,
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 509fc2f3f5b..ead412e3f6f 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -9145,7 +9145,8 @@  gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 		}
 	    }
 	  else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT
-		   && code == OMP_TARGET_EXIT_DATA)
+		   && (code == OMP_TARGET_EXIT_DATA
+		       || code == OACC_EXIT_DATA))
 	    remove = true;
 	  else if (DECL_SIZE (decl)
 		   && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
@@ -11001,8 +11002,9 @@  gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
 	   && omp_find_clause (OMP_STANDALONE_CLAUSES (expr),
 			       OMP_CLAUSE_FINALIZE))
     {
-      /* Use GOMP_MAP_DELETE/GOMP_MAP_FORCE_FROM to denote that "finalize"
-	 semantics apply to all mappings of this OpenACC directive.  */
+      /* Use GOMP_MAP_DELETE, GOMP_MAP_FORCE_DETACH, and
+	 GOMP_MAP_FORCE_FROM to denote that "finalize" semantics apply
+	 to all mappings of this OpenACC directive.  */
       bool finalize_marked = false;
       for (tree c = OMP_STANDALONE_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN (c))
 	if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
@@ -11016,6 +11018,10 @@  gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
 	      OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DELETE);
 	      finalize_marked = true;
 	      break;
+	    case GOMP_MAP_DETACH:
+	      OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_DETACH);
+	      finalize_marked = true;
+	      break;
 	    default:
 	      /* Check consistency: libgomp relies on the very first data
 		 mapping clause being marked, so make sure we did that before
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index bbcbc121bae..f5ee117887f 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -7581,6 +7581,9 @@  lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	  case GOMP_MAP_FORCE_DEVICEPTR:
 	  case GOMP_MAP_DEVICE_RESIDENT:
 	  case GOMP_MAP_LINK:
+	  case GOMP_MAP_ATTACH:
+	  case GOMP_MAP_DETACH:
+	  case GOMP_MAP_FORCE_DETACH:
 	    gcc_assert (is_gimple_omp_oacc (stmt));
 	    break;
 	  default:
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 81414d117e1..0b96011e72a 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -755,6 +755,15 @@  dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
 	case GOMP_MAP_LINK:
 	  pp_string (pp, "link");
 	  break;
+	case GOMP_MAP_ATTACH:
+	  pp_string (pp, "attach");
+	  break;
+	case GOMP_MAP_DETACH:
+	  pp_string (pp, "detach");
+	  break;
+	case GOMP_MAP_FORCE_DETACH:
+	  pp_string (pp, "force_detach");
+	  break;
 	default:
 	  gcc_unreachable ();
 	}
diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index ccfb657c735..0b9548eae3e 100644
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -42,6 +42,7 @@ 
 #define GOMP_MAP_FLAG_SPECIAL_2		(1 << 4)
 #define GOMP_MAP_FLAG_SPECIAL		(GOMP_MAP_FLAG_SPECIAL_1 \
 					 | GOMP_MAP_FLAG_SPECIAL_0)
+#define GOMP_MAP_DEEP_COPY		(1 << 5)
 /* Flag to force a specific behavior (or else, trigger a run-time error).  */
 #define GOMP_MAP_FLAG_FORCE		(1 << 7)
 
@@ -128,6 +129,13 @@  enum gomp_map_kind
     /* Decrement usage count and deallocate if zero.  */
     GOMP_MAP_RELEASE =			(GOMP_MAP_FLAG_SPECIAL_2
 					 | GOMP_MAP_DELETE),
+    /* In OpenACC, attach a pointer to a mapped struct field.  */
+    GOMP_MAP_ATTACH =			(GOMP_MAP_DEEP_COPY | 0),
+    /* In OpenACC, detach a pointer to a mapped struct field.  */
+    GOMP_MAP_DETACH =			(GOMP_MAP_DEEP_COPY | 1),
+    /* In OpenACC, detach a pointer to a mapped struct field.  */
+    GOMP_MAP_FORCE_DETACH =		(GOMP_MAP_DEEP_COPY
+					 | GOMP_MAP_FLAG_FORCE | 1),
 
     /* Internal to GCC, not used in libgomp.  */
     /* Do not map, but pointer assign a pointer instead.  */

--------------2.17.1--