[OpenACC] Implement present clause

Author: jdenny-ornl

clang/README-OpenACC-design.md

@@ -577,6 +577,7 @@ clarify these points in future versions of the OpenACC specification.
     * Data mapping attributes (DMAs), which describe the mapping and
       transfer of data between host and device:
         * `nomap` (default)
+        * `present`
         * `copy`
         * `copyin`
         * `copyout`
@@ -703,7 +704,9 @@ clarify these points in future versions of the OpenACC specification.
 * *exp*|*imp* DA for a variable of incomplete type is an error.
   Notes:
     * A private or device copy is assumed to be allocatable in each of
-      these cases, but allocation is impossible for incomplete types.
+      these cases, but creating a new allocation or checking for a
+      full existing allocation is impossible for incomplete types
+      because the size is unknown.
     * It does not appear possible for any DA other than `copy`,
       `nomap`, or `shared` to be *imp* for a variable of incomplete
       type.
@@ -741,6 +744,9 @@ clarify these points in future versions of the OpenACC specification.
     * *exp* `copyin` or *exp* `firstprivate` is fine for a `const`
       variable.  The local copy will have the original variable's
       value throughout its lifetime.
+    * *exp* `present` is fine for a `const` variable as the
+      initialization could have happened at the time of the
+      allocation.
     * *imp* `copy` or *imp* `firstprivate` for a `const` variable
       should be fine for the same reasons as their *exp* versions.
     * *imp* `nomap` and *imp* `shared` are the only remaining *imp*
@@ -1103,6 +1109,52 @@ Clacc's current mapping of an `acc data` directive and its clauses to
 OpenMP is as follows:
 
 * `acc data` -> `omp target data`
+* *exp* `present` is translated according to the
+  `-fopenacc-present-omp=KIND` command-line option:
+    * `KIND` is one of:
+        * `present` (default):
+            * *exp* `present` -> *exp* `map` with a `present,alloc`
+              map type.
+        * `alloc`:
+            * *exp* `present` -> *exp* `map` with an `alloc` map type.
+    * Notes:
+        * See the discussion of the `present` clause under "Supported
+          Features" in `README-OpenACC-status.md` for a description of
+          associated diagnostics and for an explanation of the impact
+          of this design on Clacc users.
+        * Clacc does not currently support translating `present` to
+          calls to `omp_target_is_present` for the following reasons:
+            * `omp_target_is_present` is fine for checking that data
+              is present when the data is a single byte.  However, it
+              does not provide a means to check a range of *N* bytes
+              without requiring O(*N*) calls.  Checking for just the
+              first and last bytes would be O(1), but that would not
+              handle some cases where the first and last byte are part
+              of multiple, separate, possibly non-contiguous
+              subarrays.
+            * A runtime error must be triggered when data is not
+              present.  However, we are not aware of any OpenMP 5.0
+              API for triggering a runtime error.  `fprintf`,
+              `stderr`, and `abort` might provide an obvious
+              mechanism, but that might not be how other runtime
+              errors are handled by every OpenMP runtime that might be
+              used by every OpenMP compiler.  Do they print to
+              `stderr`, a log, or both?  Does the runtime need to be
+              shut down somehow before `abort`?  Moreover, there are
+              plans to introduce a runtime-error handling extension to
+              the OpenACC Profiling Interface, but there's no obvious
+              way for it to intercept calls to `fprintf` and `abort`.
+            * Clacc would need to insert either includes of headers
+              like `stdio.h` and `omp.h` or at least local
+              declarations of symbols like `omp_target_is_present`,
+              `omp_get_default_device`, `fprintf`, `stderr`, and
+              `abort`.  Either way, such symbols could potentially
+              conflict with user-defined symbols or preprocessor macro
+              definitions.
+            * We will reconsider adding support for such a translation
+              if there is user demand, but we suspect the likelihood
+              that OpenMP TR8's `present` map type modifier will be
+              standardized means this option is likely not worthwhile.
 * *exp* `copy` -> *exp* `map` with a `tofrom` map type.
 * *exp* `copyin` -> *exp* `map` with a `to` map type.
 * *exp* `copyout` -> *exp* `map` with a `from` map type.
@@ -1222,9 +1274,8 @@ to OpenMP is as follows:
       However, this behavior shouldn't be observable given that it's
       OpenACC's structured reference counter, which is guaranteed not
       to fall to zero before the enclosing `acc data` ends either way.
-* *exp*|*imp* `copy` -> *exp* `map` with a `tofrom` map type
-* *exp* `copyin` -> *exp* `map` with a `to` map type
-* *exp* `copyout` -> *exp* `map` with a `from` map type
+* All DMAs are mapped in the same manner as when appearing on an `acc
+  data`.  *imp* `copy` is mapped in the same manner as *exp* `copy`.
 * *imp* `shared` -> *exp* `shared`
 * *exp*|*imp* `reduction` -> *exp* `reduction`
 * *exp*|*imp* `firstprivate` -> *exp* `firstprivate`

clang/README-OpenACC-status.md

@@ -20,9 +20,14 @@ We have implemented and tested support for the following features:
           tested offloading target triples.
         * See the section "Interaction with OpenMP Support" in
           `README-OpenACC-design.md` for design details.
+    * `-fopenacc-present-omp=KIND` where `KIND` is either `present` or
+      `alloc`
+        * See the discussion of the `present` clause below.
     * `-Wsource-uses-openacc`
     * `-Wopenacc-ignored-clause`
         * See the discussion of the `vector_length` clause below.
+    * `-Wopenacc-omp-map-present`
+        * See the discussion of the `present` clause below.
     * Notes:
         * See the section "Using" in `../README.md` for an
           introduction to Clacc's command-line options.
@@ -32,20 +37,79 @@ We have implemented and tested support for the following features:
 * run-time environment variables:
     * `OMP_TARGET_OFFLOAD=disabled` for targeting host
 * `data` directive:
+    * `present` clause
+        * Traditional compilation mode:
+            * The `present` clause is fully supported.
+            * Although the traditional compilation mode user typically
+              does not need to be aware, the OpenMP translation of the
+              `present` clause uses an OpenMP TR8 feature: the
+              `present` map type modifier.  This modifier is combined
+              with the OpenMP `alloc` map type as opposed to, for
+              example, `tofrom` in order to suppress device-to-host
+              transfers in the case that there is a reference count
+              decrement within the associated region.
+            * If desired, it is possible to adjust the translation or
+              related diagnostics by using the command-line options
+              discussed below for source-to-source mode.  The
+              difference is that, by default in traditional
+              compilation mode, the related diagnostics are disabled.
+        * Source-to-source mode:
+            * Occurrences of the `present` clause produce a
+              compile-time error diagnostic by default.  The purpose
+              of the diagnostic is to ensure the user is aware that
+              the translation includes the `present` map type modifier
+              because it is unlikely to be supported yet by foreign
+              OpenMP compilers.
+            * The diagnostic is actually the warning enabled by
+              `-Wopenacc-omp-map-present`, which is enabled and
+              treated as an error by default in source-to-source mode.
+            * To work around this issue, the following command-line
+              options can be specified:
+                * `-Wno-error=openacc-omp-map-present` converts the
+                  diagnostic to a warning to make it easier to find
+                  all occurrences.
+                * `-Wno-openacc-omp-map-present` disables the
+                  diagnostic entirely.  This is useful if the
+                  generated OpenMP will not be compiled or if the
+                  OpenMP compiler actually already supports the
+                  `present` map type modifier.
+                * `-fopenacc-present-omp=alloc` suppresses the
+                  diagnostic by changing the translation to use the
+                  standard OpenMP `alloc` map type without the
+                  `present` modifier.  In this case, there is no
+                  runtime error when the specified variable is not
+                  present on the device.  However, this translation
+                  should be sufficient for OpenACC applications that
+                  are robust enough not to actually encounter this
+                  runtime error.
+        * `-fopenacc[-ast]-print=acc`, `-ast-print`, `-ast-dump`,
+          etc. mode:
+            * Debugging modes like these do not actually print OpenMP
+              source code, so they leave the aforementioned diagnostic
+              disabled as in traditional compilation mode.
+        * Currently, Clacc's implementation of the OpenMP TR8
+          `present` map type modifier is not well tested outside of
+          Clacc translations from OpenACC to OpenMP.  Thus, it is not
+          yet recommended for use in hand-written OpenMP code as it
+          might not integrate well with some OpenMP features.
+        * See the section "Data Directives" in
+          `README-OpenACC-design.md` for a discussion of alternative
+          translations that we considered for the OpenACC `present`
+          clause.
     * `copy` clause and aliases `pcopy` and `present_or_copy`
     * `copyin` clause and aliases `pcopyin` and
       `present_or_copyin`
     * `copyout` clause and aliases `pcopyout` and
       `present_or_copyout`
-    * in `copy`, `copyin`, and `copyout` clauses and their aliases,
-      subarrays specifying contiguous blocks
+    * in `present`, `copy`, `copyin`, and `copyout` clauses and their
+      aliases, subarrays specifying contiguous blocks
     * any number of levels of nesting within other `data` directives
 * `parallel` directive:
     * use without clauses
     * data attributes:
         * implicit `copy` for non-scalars
         * implicit `firstprivate` for scalars
-        * For `copy`, `copyin`, and `copyout` clauses and
+        * For `present`, `copy`, `copyin`, and `copyout` clauses and
           their aliases, support is the same as for the `data`
           directive, as described above.
         * `firstprivate` clause
@@ -288,3 +352,26 @@ Clacc matures.  Please report any cases not listed below.
       class to something more general like `ArrayRangeExpr` or
       `ArraySubscriptExtendedExpr`.  See the todo on this class in the
       implementation.
+* OpenMP runtime error diagnostics are expressed in terms of OpenMP
+  not OpenACC.
+
+OpenMP Extensions
+=================
+
+There are some OpenACC features for which Clacc depends on OpenMP
+extensions because we are not aware of standard OpenMP features that
+are sufficient.  While this dependence should not affect traditional
+compilation using Clacc's compiler and OpenACC/OpenMP runtime, it can
+affect compilation using Clacc's source-to-source mode followed by a
+foreign OpenMP compiler or runtime.  A goal of Clacc is to rely on
+standard OpenMP as much as possible, and to that end it might be
+worthwhile to propose these extensions for inclusion in the OpenMP
+specification.  Currently, Clacc uses OpenMP extensions as follows:
+
+* The `present` clause translation depends on the OpenMP TR8 map type
+  modifier `present` by default:
+    * See the discussion of the `present` clause above for details.
+* Some features of the OpenACC Profiling Interface depend on OMPT
+  extensions:
+    * See the section "OpenACC Profiling Interface" in
+      `README-OpenACC-design.md` for details.

clang/include/clang/AST/OpenACCClause.h

@@ -87,6 +87,8 @@ class ACCClause {
   SourceLocation getBeginLoc() const { return StartLoc; }
   /// Returns the ending location of the clause.
   SourceLocation getEndLoc() const { return EndLoc; }
+  /// Returns the source range of the clause.
+  SourceRange getSourceRange() const { return SourceRange(StartLoc, EndLoc); }
 
   typedef StmtIterator child_iterator;
   typedef ConstStmtIterator const_child_iterator;
@@ -244,6 +246,70 @@ class ACCNomapClause final
   }
 };
 
+/// This represents the clause 'present' (or any of its aliases) for
+/// '#pragma acc ...' directives.
+///
+/// \code
+/// #pragma acc parallel present(a,b)
+/// \endcode
+/// In this example directive '#pragma acc parallel' has clause 'present' with
+/// the variables 'a' and 'b'.
+class ACCPresentClause final
+    : public ACCVarListClause<ACCPresentClause>,
+      private llvm::TrailingObjects<ACCPresentClause, Expr *> {
+  friend TrailingObjects;
+  friend ACCVarListClause;
+  friend class ACCClauseReader;
+
+  /// Build clause with number of variables \a N.
+  ///
+  /// \param Determination How the clause was determined.
+  /// \param StartLoc Starting location of the clause.
+  /// \param LParenLoc Location of '('.
+  /// \param EndLoc Ending location of the clause.
+  /// \param N Number of the variables in the clause.
+  ACCPresentClause(OpenACCDetermination Determination, SourceLocation StartLoc,
+                   SourceLocation LParenLoc, SourceLocation EndLoc, unsigned N)
+      : ACCVarListClause<ACCPresentClause>(ACCC_present, Determination,
+                                           StartLoc, LParenLoc, EndLoc, N)
+  {}
+
+  /// Build an empty clause.
+  ///
+  /// \param N Number of variables.
+  explicit ACCPresentClause(unsigned N)
+      : ACCVarListClause<ACCPresentClause>(ACCC_present, N) {
+  }
+
+public:
+  /// Creates clause with a list of variables \a VL.
+  ///
+  /// \param C AST context.
+  /// \param Determination How the clause was determined.
+  /// \param StartLoc Starting location of the clause.
+  /// \param LParenLoc Location of '('.
+  /// \param EndLoc Ending location of the clause.
+  /// \param VL List of references to the variables.
+  static ACCPresentClause *Create(
+      const ASTContext &C, OpenACCDetermination Determination,
+      SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation EndLoc,
+      ArrayRef<Expr *> VL);
+  /// Creates an empty clause with the place for \a N variables.
+  ///
+  /// \param C AST context.
+  /// \param N The number of variables.
+  static ACCPresentClause *CreateEmpty(const ASTContext &C, unsigned N);
+
+  child_range children() {
+    return child_range(reinterpret_cast<Stmt **>(varlist_begin()),
+                       reinterpret_cast<Stmt **>(varlist_end()));
+  }
+
+  static bool classof(const ACCClause *T) {
+    return T->getClauseKind() == ACCC_present;
+  }
+};
+
 /// This represents the clause 'copy' (or any of its aliases) for
 /// '#pragma acc ...' directives.
 ///

clang/include/clang/AST/RecursiveASTVisitor.h

@@ -3622,6 +3622,12 @@ bool RecursiveASTVisitor<Derived>::VisitACCNomapClause(ACCNomapClause *C) {
   return true;
 }
 
+template <typename Derived>
+bool RecursiveASTVisitor<Derived>::VisitACCPresentClause(ACCPresentClause *C) {
+  TRY_TO(VisitACCClauseList(C));
+  return true;
+}
+
 template <typename Derived>
 bool RecursiveASTVisitor<Derived>::VisitACCCopyClause(ACCCopyClause *C) {
   TRY_TO(VisitACCClauseList(C));

clang/include/clang/Basic/DiagnosticGroups.td

@@ -1113,6 +1113,7 @@ def OpenMP : DiagGroup<"openmp", [
 // OpenACC warnings.
 def SourceUsesOpenACC : DiagGroup<"source-uses-openacc">;
 def OpenACCIgnoredClause : DiagGroup<"openacc-ignored-clause">;
+def OpenACCOMPMapPresent : DiagGroup<"openacc-omp-map-present">;
 
 // Backend warnings.
 def BackendInlineAsm : DiagGroup<"inline-asm">;

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -10368,6 +10368,17 @@ def note_acc_implied_as_gang_reduction : Note<
   "implied as gang reduction here">;
 def err_acc_no_data_clause : Error<
   "expected at least one data clause for '#pragma acc %0'">;
+def warn_acc_omp_map_present : Warning<
+  "the OpenACC '%0' clause translation uses the 'present' map type modifier, "
+  "which is an OpenMP TR8 feature that might not be supported yet by other "
+  "compilers">,
+  InGroup<OpenACCOMPMapPresent>, DefaultIgnore;
+def note_acc_alternate_omp : Note<
+  "an alternative OpenMP translation can be specified with, for example, "
+  "'-fopenacc-%0-omp=%1'">;
+def note_acc_disable_diag : Note<
+  "or you can just disable this diagnostic with "
+  "'-Wno-openacc-omp-map-present'">;
 } // end of OpenACC category
 
 let CategoryName = "Related Result Type Issue" in {

clang/include/clang/Basic/LangOptions.def

@@ -231,6 +231,9 @@ LANGOPT(OpenMPCUDABlocksPerSM  , 32, 0, "Number of blocks per SM for CUDA device
 LANGOPT(OpenMPCUDAReductionBufNum , 32, 1024, "Number of the reduction records in the intermediate reduction buffer used for the teams reductions.")
 LANGOPT(OpenMPOptimisticCollapse  , 1, 0, "Use at most 32 bits to represent the collapsed loop nest counter.")
 LANGOPT(OpenACC           , 1, 0, "OpenACC support")
+ENUM_LANGOPT(OpenACCPresentOMP, OpenACCPresentOMPKind, 1,
+             OpenACCPresentOMP_Present,
+             "The OpenMP translation of the OpenACC 'present' clause")
 LANGOPT(RenderScript      , 1, 0, "RenderScript")
 
 LANGOPT(CUDAIsDevice      , 1, 0, "compiling for CUDA device")

clang/include/clang/Basic/LangOptions.h

@@ -108,6 +108,21 @@ class LangOptions : public LangOptionsBase {
     DCC_RegCall
   };
 
+  enum OpenACCPresentOMPKind {
+    OpenACCPresentOMP_Present,
+    OpenACCPresentOMP_Alloc,
+    OpenACCPresentOMP_Last = OpenACCPresentOMP_Alloc
+  };
+  static StringRef getOpenACCPresentOMPValue(OpenACCPresentOMPKind K) {
+    switch (K) {
+    case OpenACCPresentOMP_Present:
+      return "present";
+    case OpenACCPresentOMP_Alloc:
+      return "alloc";
+    }
+    llvm_unreachable("unexpected OpenACCPresentOMPKind");
+  }
+
   enum AddrSpaceMapMangling { ASMM_Target, ASMM_On, ASMM_Off };
 
   // Corresponds to _MSC_VER