Commit 25f67dfc authored by Joel E. Denny's avatar Joel E. Denny
Browse files

[Clacc][OpenACC] Implement async clause on acc parallel

See the status document for current limitations.
parent 1e8f734a
Loading
Loading
Loading
Loading
Loading
+33 −0
Original line number Diff line number Diff line
@@ -1783,6 +1783,39 @@ to OpenMP is as follows:
    * A constant expression argument here might be used by a nested
      vector-partitioned `acc loop`, but a non-constant-expression
      argument is not (this follows "Semantic Clarifications" above).
* *exp* `async` with no argument ->
  `nowait depend(inout:*acc2omp_async2dep(acc_async_noval))`
* *exp* `async(`*async_arg*`)` such that *async_arg* is a constant integer
  expression whose value is `acc_async_sync` -> discarded by the translation
* *exp* `async(`*async_arg*`)` otherwise ->
  `nowait depend(inout:*acc2omp_async2dep(`*async_arg*`))`
* Wrap the `omp target teams` and its associated statement in a compound
  statement and insert an
  `omp taskwait depend(inout:*acc2omp_async2dep(acc_async_sync))` directive at
  the end of it if any of the following conditions hold:
    * *exp* `async` with no argument.
    * *exp* `async(`*async_arg*`)` such that *async_arg* is a constant integer
      expression whose value is `acc_async_noval`.
    * *exp* `async(`*async_arg*`)` such that *async_arg* is a non-constant
      expression of signed type.
* Notes for the `async` translation:
    * `char *acc2omp_async2dep(int)` is prototyped in Clacc's `openacc.h` and
      implemented in libacc2omp.  It maps each *async_arg* to an lvalue
      appropriate for an OpenMP `depends` clauses.  See the discussion of it in
      `README-OpenACC-status.md`.
    * As described above, `nowait` and `depend` are needed if the selected
      activity queue might be an asynchronous activity queue at run time.  When
      that possibility exists but the possibility also exists that the selected
      activity queue might be the synchronous activity queue, then the trailing
      `omp taskwait` is also needed.
    * When *async_arg* is a constant integer expression whose value is
      `acc_async_default`, Clacc's Clang assumes the selected activity queue is
      asynchronous because it assumes libacc2omp's implementation, which selects
      queue 0 as the initial default activity queue.
    * The enclosing compound statement in the case of the `omp taskwait` is not
      strictly necessary in many contexts.  In the Clang AST, it simply makes it
      easier to attach the full OpenMP translation to the original OpenACC node
      because the compound statement subtree has a single root node.

Loop Directives
---------------
+43 −3
Original line number Diff line number Diff line
@@ -118,9 +118,9 @@ OpenACC-related and OpenMP-related command-line options, run Clacc's
        * Each option will be removed when Clacc develops full support for the
          associated features.
    * `-fopenacc-fake-async-wait`
        * Clacc accepts but discards OpenACC directives and clauses associated
          with async/wait support.  That is, they have no OpenMP translation in
          source-to-source mode.
        * Clacc accepts but discards some OpenACC directives and clauses
          associated with async/wait support.  That is, they have no OpenMP
          translation in source-to-source mode.
        * Clacc inserts preprocessor definitions to handle OpenACC Runtime
          Library API routines and other symbols associated with async/wait
          support.
@@ -128,6 +128,19 @@ OpenACC-related and OpenMP-related command-line options, run Clacc's
        * Some async/wait features might not be covered yet and thus will
          still produce compile-time diagnostics.  We are adding them as the
          need arises in the applications we are investigating.
        * Some async/wait features are starting to be supported when
          `-fopenacc-fake-async-wait` is not specified.  They are documented
          along with other supported features below.  However,
          `-fopenacc-fake-async-wait` fakes even those async features that
          Clacc otherwise supports because they are not safe while it fakes at
          least some wait features.  Because that produces synchronous behavior,
          wait features have no effect on correctness, so
          `-fopenacc-fake-async-wait` does not bother to fake wait features that
          are already supported.
        * Summary: If an application uses only fully supported async/wait
          features, it doesn't need `-fopenacc-fake-async-wait`.  Otherwise,
          `-fopenacc-fake-async-wait` might enable the application to compile
          and run successfully, but it will eliminate the asynchronous behavior.
    * `-fopenacc-fake-tile-clause`
        * Clacc now fully supports the `tile` clause, so this option is
          deprecated and has no effect.
@@ -319,6 +332,27 @@ Run-Time Environment Variables
          `simdlen`.  OpenACC does permit the compiler to ignore
          `vector_length` as a hint, so we choose to ignore it and
          warn in the case of a non-constant expression.
* `async` clause
    * The argument can be omitted, or it must be an integer expression.  It must
      evaluate to a non-negative integer or the value of `acc_async_sync`,
      `acc_async_noval`, or `acc_async_default` as defined in Clacc's
      `openacc.h`.  An argument that is not of integer type or that is an
      integer constant expression that does not meet the above restrictions
      produces a compile-time error diagnostic.  Otherwise, an expression that
      does not meet the above restrictions produces a runtime error.
    * The function prototype `char *acc2omp_async2dep(int)` must be in scope
      and must link to Clacc's implementation of it in `libacc2omp.so`.  The
      easiest way to accomplish this is usually to just add
      `#include <openacc.h>`.  In the future, Clacc might insert the prototype
      automatically where it is not in scope.  Also see "Source-to-Source Mode
      Limitations" below.
    * Activity queues are currently common to all offload devices instead of per
      device, limiting some concurrency.  In the future, this limitation might
      be removed.
    * Due to a nondeterminism apparently inherited from upstream Clang's OpenMP
      implementation, concurrent execution of activity queues is not guaranteed.
      Moreover, there are occasional runtime assertion failures when targeting
      the host.  We need to investigate these issues further.

`loop` Directive
----------------
@@ -982,6 +1016,9 @@ Language Support
Source-to-Source Mode Limitations
=================================

These limitations affect source-to-source mode but have no effect on traditional
compilation mode.

* Calls to the OpenACC Runtime Library API are not translated to
  OpenMP at compile time:
    * `acc_on_device` is implemented in terms of OpenMP fully within
@@ -994,6 +1031,9 @@ Source-to-Source Mode Limitations
      Runtime Library Routines and require Clacc's OpenACC runtime
      library to be linked.  See the section "Linking" in
      `../README.md` for details.
* Calls to `acc2omp_async2dep` are used in the translation of `async` clauses
  but are not specified by OpenMP.  It also requires Clacc's OpenACC runtime
  library to be linked.
* Occurrences of the `_OPENACC` preprocessor macro are not translated
  to OpenMP:
    * Instead, the compiler inserts its `_OPENACC` definition at the
+58 −5
Original line number Diff line number Diff line
@@ -1836,30 +1836,76 @@ public:
class ACCAsyncClause : public ACCClause {
  friend class ACCClauseReader;

public:
  enum {
    // These enumerator values must be kept in sync with
    // openmp/libacc2omp/src/include/openacc.h.var.
    Acc2ompAccAsyncSync = -1,
    Acc2ompAccAsyncNoval = -2,
    Acc2ompAccAsyncDefault = -3,
  };

  enum AsyncArgStatus {
    AsyncArgIsSync, ///< definitely the synchronous activity queue
    AsyncArgIsAsync, ///< definitely an asynchronous activity queue
    AsyncArgIsUnknown, ///< might be synchronous or asynchronous activity queue
    AsyncArgIsError, ///< invalid async-arg
  };

  static constexpr StringRef Async2DepName = "acc2omp_async2dep";

private:
  /// Location of '('.
  SourceLocation LParenLoc;

  /// Original AsyncArg expression.
  Stmt *AsyncArg = nullptr;

  /// Status of AsyncArg.
  AsyncArgStatus TheAsyncArgStatus = AsyncArgIsError;

  /// The \c acc2omp_async2dep symbol found in scope at this clause, or
  /// \c nullptr if none.
  NamedDecl *Async2Dep = nullptr;

  /// Set the original AsyncArg expression.
  ///
  /// \param E AsyncArg expression.
  void setAsyncArg(Expr *E) { AsyncArg = E; }

  /// Set AsyncArg status, which must not be \c AsyncArgIsError.
  void setAsyncArgStatus(AsyncArgStatus TheAsyncArgStatus) {
    assert(TheAsyncArgStatus != AsyncArgIsError &&
           "expected valid async-arg status");
    this->TheAsyncArgStatus = TheAsyncArgStatus;
  }

  /// Set the \c acc2omp_async2dep symbol found in scope at this clause, or
  /// \c nullptr if none.
  void setAsync2Dep(NamedDecl *ND) { Async2Dep = ND; }

public:
  /// Build 'async' clause.
  ///
  /// \param E Original expression associated with this clause, or \c nullptr if
  ///        omitted.
  /// \param AsyncArg Original expression associated with this clause, or
  ///        \c nullptr if omitted.
  /// \param TheAsyncArgStatus Status of \c AsyncArg.  Must not be
  ///        \c AsyncArgIsError.
  /// \param Async2Dep The \c acc2omp_async2dep symbol found in scope at this
  ///        clause, or \c nullptr if none.
  /// \param StartLoc Starting location of the clause.
  /// \param LParenLoc Location of '(', or an invalid location if argument
  ///        omitted.
  /// \param EndLoc Ending location of the clause.
  ACCAsyncClause(Expr *E, SourceLocation StartLoc, SourceLocation LParenLoc,
                 SourceLocation EndLoc)
  ACCAsyncClause(Expr *AsyncArg, AsyncArgStatus TheAsyncArgStatus,
                 NamedDecl *Async2Dep, SourceLocation StartLoc,
                 SourceLocation LParenLoc, SourceLocation EndLoc)
      : ACCClause(ACCC_async, ACC_EXPLICIT, StartLoc, EndLoc),
        LParenLoc(LParenLoc), AsyncArg(E) {}
        LParenLoc(LParenLoc), AsyncArg(AsyncArg),
        TheAsyncArgStatus(TheAsyncArgStatus), Async2Dep(Async2Dep) {
    assert(TheAsyncArgStatus != AsyncArgIsError &&
           "expected valid async-arg status");
  }

  /// Build an empty clause.
  ACCAsyncClause() : ACCClause(ACCC_async) {}
@@ -1876,6 +1922,13 @@ public:
  /// Return the original AsyncArg expression or \c nullptr if omitted.
  Expr *getAsyncArg() const { return cast_or_null<Expr>(AsyncArg); }

  /// Return AsyncArg status, which is never \c AsyncArgIsError.
  AsyncArgStatus getAsyncArgStatus() const { return TheAsyncArgStatus; }

  /// Get the \c acc2omp_async2dep symbol found in scope at this clause, or
  /// \c nullptr if none.
  NamedDecl *getAsync2Dep() const { return Async2Dep; }

  child_range children() { return child_range(&AsyncArg, &AsyncArg + 1); }

  static bool classof(const ACCClause *T) {
+7 −0
Original line number Diff line number Diff line
@@ -11362,6 +11362,13 @@ def note_acc_routine_implicit : Note<
def err_acc_routine_for_orphaned_loop : Error<
  "function '%0' has no explicit '#pragma acc routine' but contains orphaned "
  "'#pragma acc loop'">;
def err_acc_missing_function : Error<"function '%0' prototype is not in scope">;
def err_acc_not_function : Error<"'%0' is not a function of type %1">;
def note_acc_required_by_clause : Note <"required by '%0' clause here">;
def note_acc_include_openacc_h : Note<"try including openacc.h">;
def err_acc_clause_not_async_arg : Error<
  "argument to '%0' clause must be a non-negative integer, acc_async_sync, "
  "acc_async_noval, or acc_async_default">;
} // end of OpenACC category
let CategoryName = "Related Result Type Issue" in {
+9 −5
Original line number Diff line number Diff line
@@ -3531,7 +3531,6 @@ public:
  ///        directive.  It's indexed by \c OpenACCClauseKind to record which
  ///        clauses have been seen except it omits those that are never
  ///        permitted on the directive.
  ///
  ACCClause *ParseOpenACCClause(OpenACCDirectiveKind DKind,
                                OpenACCClauseKind CKind,
                                SmallVectorImpl<bool> &SeenClauses);
@@ -3540,9 +3539,11 @@ public:
  /// \param Kind Kind of current clause.
  /// \param ParseOnly true to skip the clause's semantic actions and return
  /// nullptr.
  ///
  /// \param Async2Dep The currently visible declaration for
  ///        \c acc2omp_async2dep, or \c nullptr if none or if \c CKind is not
  ///        \c ACCC_async.
  ACCClause *ParseOpenACCSingleExprClause(OpenACCClauseKind Kind,
                                          bool ParseOnly);
                                          bool ParseOnly, NamedDecl *Async2Dep);
  /// Parses OpenACC 'gang' clause with an argument.
  ///
  /// \param ParseOnly true to skip the clause's semantic actions and return
@@ -3554,8 +3555,11 @@ public:
  /// \param Kind Kind of current clause.
  /// \param ParseOnly true to skip the clause's semantic actions and return
  /// nullptr.
  ///
  ACCClause *ParseOpenACCNoArgClause(OpenACCClauseKind Kind, bool ParseOnly);
  /// \param Async2Dep The currently visible declaration for
  ///        \c acc2omp_async2dep, or \c nullptr if none or if \c CKind is not
  ///        \c ACCC_async.
  ACCClause *ParseOpenACCNoArgClause(OpenACCClauseKind Kind, bool ParseOnly,
                                     NamedDecl *Async2Dep);
  /// Parses clause with the list of variables of a kind \a Kind.
  ///
  /// \param Kind Kind of current clause.
Loading