[Clacc][OpenACC] num_workers -> thread_limit

* *exp*|*imp* `firstprivate` -> *exp* `firstprivate`

* *exp* `private` -> *exp* `private`

* *exp* `num_gangs` -> *exp* `num_teams`

* If *exp* `num_workers` with a non-constant-expression argument, and

  if there is a nested worker-partitioned `acc loop`, then *exp*

  `num_workers` -> wrap the `omp target teams` in a compound statement

  and declare a local `const` variable with the same type and value as

  the *exp* `num_workers` argument.

* Else if *exp* `num_workers` with a non-constant-expression argument

  that potentially has side effects, then *exp* `num_workers` -> wrap

  the `omp target teams` in a compound statement and insert a

  statement that casts the argument's expression to `void`.

* Else, translation discards *exp* `num_workers`.  Notes:

    * A constant-expression argument here might be used by a nested

      worker-partitioned `acc loop`.

* *exp* `num_workers` -> *exp* `thread_limit`

* If *exp* `vector_length` with a non-constant-expression argument

  that potentially has side effects, then *exp* `vector_length` ->

  wrap the `omp target teams` in a compound statement and insert a

* The output `distribute`, `parallel for`, and `simd` OpenMP directive

  components are sorted in the above order before all clauses regardless of the

  input clause order.

* If *exp*|*imp* `worker` and either *not* `tile` or *exp*|*imp* `vector`, then

  *exp* `num_workers` from ancestor `acc parallel` -> *exp* `num_threads` where

  the argument is either (1) the original *exp* `num_workers` argument if it is

  a constant expression or (2) otherwise an expression containing only a

  reference to the local `const` variable generated for that *exp*

  `num_workers`.  Notes:

    * For the ancestor `acc parallel` and for all OpenACC directives

      nested between it and this `acc loop`, Clacc leaves the OpenMP

      data sharing attribute for the local `const` variable for

      `num_workers` as implicit.  Because the variable is `const`,

      private copies are not useful, so sharing is probably most

      efficient, but not all OpenMP directives permit an *exp*

      `shared` clause.  Thus, relying on implicit data sharing

      attributes throughout simplifies the implementation.

* If *exp*|*imp* `vector` and *not* `tile`, then *exp* `vector_length` with a

  constant-expression argument from ancestor `acc parallel` -> *exp* `simdlen`.

* `static:*` within `gang` -> `dist_schedule(static)`

  computed automatically.  If `acc loop vector` were mapped to `omp

  parallel for`, `vector_length` with a non-constant-expression

  argument would be possible.

* Orphaned `acc loop` directive that observes `num_workers` and

  `vector_length` because the enclosing compute construct from which

  those clauses would normally be applied during translation is not

  statically visible.

* Orphaned `acc loop` directive that observes `vector_length` because the

  enclosing compute construct from which `vector_length` would normally be

  copied during translation is not statically visible.

* Subarrays specifying non-contiguous blocks in dynamic

  multidimensional arrays because these cannot be mapped to OpenMP

  array sections.  Notes:

            * `num_gangs`, `num_workers`, and `vector_length` are

              omitted because we do not know how to obtain them via

              OMPT:

                * The problem with `num_gangs` is that OpenACC 2.7

                  says it's the number of gangs *created*, but the

                  `ompt_callback_target_submit` callback only provides

                  the number of teams *requested*.  It might possible

                  to retrieve the required data from OMPT trace

                  records, but we have not implemented that support

                  yet.

                * The problem with `num_workers` and `vector_length`

                  is that, in contrast with OpenACC compute

                  directives, `num_threads` and `simdlen` are not

                  specified at the level of an OpenMP target

                  directive.

                * OpenACC 3.3, sec. 5.2.2, p. 132 says they're the "number of

                  gangs, workers, and vector lanes created for this kernel

                  launch."

                * The problem with `num_gangs` is that the

                  `ompt_callback_target_submit` callback provides only the

                  number of teams *requested*.  It might be possible to retrieve

                  the number of teams *created* from OMPT trace records, but we

                  have not implemented that support yet.

                * The problem with `num_workers` is that the `thread_limit`

                  clause at the `omp target teams` directive specifies only the

                  *limit* on the number of threads to be created *later* at an

                  OpenMP `parallel` directive.  There appears to be no data

                  provided by the `ompt_callback_target_submit` callback or

                  associated OMPT trace records on the actual number of threads

                  that will be created at that later point in time.  We might

                  investigate whether an OpenMP Runtime Library Routine can be

                  called to accurately predict this value.

                * The problem with `vector_length` is that `simdlen` is

                  specified later on a `simd` directive rather than at the

                  `omp target teams` directive.

    * `acc_api_info`:

        * `device_api` is always set to `acc_device_api_none` because

          it's used to indicate the semantics of later fields we do

      does not use the clause and reports a warning diagnostic, which

      can be suppressed or converted to an error using the

      `-W{no-,error=}openacc-ignored-clause` command-line options.

    * `num_workers` and `vector_length` currently do not affect

      orphaned `loop` constructs.

    * `vector_length` currently does not affect orphaned `loop` constructs.

    * Notes:

        * OpenACC 2.6 specifies only that the arguments must be

          integer expressions.  However, OpenMP specifies the stricter

          requirements above for `num_teams`, `num_threads`, and

          `simdlen`, to which Clacc translates the above clauses.

        * In the cases of `num_workers` and `vector_length`, OpenACC 3.3

          specifies only that the arguments must be integer expressions.

          However, OpenMP 5.2 specifies the stricter requirements above for

          `thread_limit` and `simdlen`, to which Clacc translates them.

        * A non-positive value here probably doesn't make sense

          anyway.  Moreover, if the argument is an integer constant

          (so that it can be statically analyzed), gcc 7.3.0 warns if

          directive" below), select each loop nest's outermost `loop` constructs

          on which `worker` and `vector` clauses are permitted.  A more

          sophisticated analysis might be employed in the future.

        * Currently, Clang might misbehave when a `loop` construct receives an

          implicit `worker` clause and appears within a `parallel` construct

          that has a `num_workers` clause with a non-constant expression.

    * For now, if none of these clauses appear (explicitly or

      implicitly), then a sequential loop is produced.

* Supported data attributes and clauses

                                              SourceLocation StartLoc,

                                              SourceLocation LParenLoc,

                                              SourceLocation EndLoc) {

  // OpenMP says num_threads must evaluate to a positive integer value.

  // OpenMP says thread_limit must evaluate to a positive integer value.

  // OpenACC doesn't specify such a restriction that I see for num_workers, but

  // it seems reasonable.

  if (PosIntError == IsPositiveIntegerValue(NumWorkers, *this, ACCC_num_workers,

    llvm::DenseMap<ACCDataVar, DAVarData> DAMap;

    ///@{

    /// Before translating the associated statement of an acc parallel

    /// directive, TransformACCParallelDirective sets these as follows, and they

    /// are copied to descendant stack entries as those entries are created.

    /// directive, TransformACCParallelDirective sets this as follows, and it

    /// is copied to descendant stack entries as those entries are created.

    ///

    /// If the acc parallel directive has a num_workers or constant-expression

    /// vector_length clause, then NumWorkersExpr or VectorLengthExpr,

    /// respectively, is set to its value.

    ///

    /// If the acc parallel directive has a num_workers clause with a

    /// non-constant expression and it has a (separate or combined with this

    /// directive) nested acc loop directive with worker partitioning, then

    /// NumWorkersVarDecl is set to the declaration of a constant variable

    /// generated for the sake of OpenMP and initialized with NumWorkersExpr.

    VarDecl *NumWorkersVarDecl = nullptr;

    Expr *NumWorkersExpr = nullptr;

    /// If the acc parallel directive has a constant-expression vector_length

    /// clause, then VectorLengthExpr is set to its value.

    Expr *VectorLengthExpr = nullptr;

    ///@}

  };

      }

      // Copy data from parent directive.

      DirEntry.NumWorkersVarDecl = ParentDirEntry.NumWorkersVarDecl;

      DirEntry.NumWorkersExpr = ParentDirEntry.NumWorkersExpr;

      DirEntry.VectorLengthExpr = ParentDirEntry.VectorLengthExpr;

    }

    /// Pop top directive's data from Transform.DirStack.

    ConditionalCompoundStmtRAII(TransformACCToOMP &Tx)

        : Tx(Tx), Started(false), Finalized(false), Err(false) {}

    /// Add a variable declaration that privatizes an enclosing declaration.

    /// Use \c addNewPrivateDecl instead if it's an entirely new variable.

    /// (For an entirely new variable, we once had an \c addNewPrivateDecl, but

    /// there are currently no uses.)

    void addPrivatizingDecl(SourceLocation RefStartLoc, SourceLocation RefEndLoc,

                            VarDecl *VD) {

      prepForAdd();

          Tx.getSema().ConvertDeclToDeclGroup(DPrivate), RefStartLoc,

          RefEndLoc));

    }

    /// Add a local variable declaration that doesn't privatize an enclosing

    /// declaration.  Use \c addPrivatizingDecl instead if it does.

    VarDecl *addNewPrivateDecl(StringRef Name, QualType Ty, Expr *Init,

                               SourceLocation Loc) {

      prepForAdd();

      ASTContext &Context = Tx.getSema().getASTContext();

      DeclContext *DC = Tx.getSema().CurContext;

      IdentifierInfo *II = &Tx.getSema().PP.getIdentifierTable().get(Name);

      TypeSourceInfo *TInfo = Context.getTrivialTypeSourceInfo(Ty, Loc);

      VarDecl *VD = VarDecl::Create(Context, DC, Loc, Loc, II, Ty, TInfo,

                                    SC_None);

      Tx.getSema().AddInitializerToDecl(VD, Init, false);

      add(Tx.getSema().ActOnDeclStmt(

          Tx.getSema().ConvertDeclToDeclGroup(VD), VD->getBeginLoc(),

          VD->getEndLoc()));

      return VD;

    }

    /// Evaluate an expression that has side effects but whose original use has

    /// been removed.

    void addUnusedExpr(Expr *E) {

    ConditionalCompoundStmtRAII EnclosingCompoundStmt(*this);

    ASTContext &Context = getSema().getASTContext();

    // Declare a num_workers variable in an enclosing compound statement, if

    // needed.  FIXME: This generates an unused __clang_acc_num_workers__

    // declaration when the only loop worker clause is discarded due to a tile

    // clause or when there's just a call to a worker function.  Moreover, it

    // isn't generated when there's a loop with an implicit worker clause.

    auto NumWorkersClauses = D->getClausesOfKind<ACCNumWorkersClause>();

    if (NumWorkersClauses.begin() != NumWorkersClauses.end()) {

      DirEntry.NumWorkersExpr = NumWorkersClauses.begin()->getNumWorkers();

      if (D->getNestedExplicitWorkerPartitioning()) {

        if (!DirEntry.NumWorkersExpr->isIntegerConstantExpr(Context))

          DirEntry.NumWorkersVarDecl = EnclosingCompoundStmt.addNewPrivateDecl(

              "__clang_acc_num_workers__",

              DirEntry.NumWorkersExpr->getType().withConst(),

              DirEntry.NumWorkersExpr, DirEntry.NumWorkersExpr->getBeginLoc());

      } else if (DirEntry.NumWorkersExpr->HasSideEffects(Context))

        EnclosingCompoundStmt.addUnusedExpr(DirEntry.NumWorkersExpr);

    }

    auto VectorLengthClauses = D->getClausesOfKind<ACCVectorLengthClause>();

    if (VectorLengthClauses.begin() != VectorLengthClauses.end()) {

      Expr *E = VectorLengthClauses.begin()->getVectorLength();

    // What kind of OpenMP directive should we build?

    // OMPD_unknown means none (so sequential).

    OpenMPDirectiveKind TDKind;

    Expr *AddNumThreadsExpr = nullptr;

    Expr *AddSimdlenExpr = nullptr;

    bool AddScopeWithLCVPrivate = false;

    bool AddScopeWithAllPrivates = false;

      TDKind = OMPD_unknown;

      AddScopeWithAllPrivates = true;

    } else {

      if (Partitioning.hasWorkerPartitioning()) {

        if (DirEntry.NumWorkersVarDecl) {

          ExprResult Res = getSema().BuildDeclRefExpr(

              DirEntry.NumWorkersVarDecl,

              DirEntry.NumWorkersVarDecl->getType().getNonReferenceType(),

              VK_PRValue, D->getEndLoc());

          assert(!Res.isInvalid() &&

                 "expected valid reference to num_workers variable");

          AddNumThreadsExpr = Res.get();

        } else

          AddNumThreadsExpr = DirEntry.NumWorkersExpr;

      }

      if (Partitioning.hasVectorPartitioning()) {

        AddSimdlenExpr = DirEntry.VectorLengthExpr;

        AddScopeWithLCVPrivate = true;

    getSema().StartOpenMPDSABlock(TDKind, DeclarationNameInfo(),

                                  /*CurScope=*/nullptr, D->getBeginLoc());

    // Add num_threads and simdlen clauses, as needed.

    // Add simdlen clause if needed.

    llvm::SmallVector<OMPClause *, 16> TClauses;

    size_t NumClausesAdded = 0;

    if (AddNumThreadsExpr) {

      OpenMPStartEndClauseRAII ClauseRAII(getSema(), OMPC_num_threads);

      TClauses.push_back(getDerived().RebuildOMPNumThreadsClause(

          AddNumThreadsExpr, AddNumThreadsExpr->getBeginLoc(),

          AddNumThreadsExpr->getBeginLoc(), AddNumThreadsExpr->getEndLoc()));

      ++NumClausesAdded;

    }

    if (AddSimdlenExpr) {

      OpenMPStartEndClauseRAII ClauseRAII(getSema(), OMPC_simdlen);

      TClauses.push_back(getDerived().RebuildOMPSimdlenClause(

  OMPClauseResult TransformACCNumWorkersClause(ACCDirectiveStmt *D,

                                               OpenMPDirectiveKind TDKind,

                                               ACCNumWorkersClause *C) {

    return OMPClauseEmpty();

    ExplicitClauseLocs L(D, C, C->getLParenLoc());

    return getDerived().RebuildOMPThreadLimitClause(

        C->getNumWorkers(), L.LocStart, L.LParenLoc, L.LocEnd);

  }

  OMPClauseResult TransformACCVectorLengthClause(ACCDirectiveStmt *D,

// DEFINE: %{check}( PRT_ARG %, PRT_CHK %, PRT_VER %) =                        \

// DEFINE:   %clang -Xclang -verify=%{PRT_VER} -fopenacc-print=%{PRT_ARG}      \

// DEFINE:          -Wno-openacc-omp-ext -ferror-limit=100 %t-acc.c |          \

// DEFINE:          -Wno-openacc-ignored-clause  -Wno-openacc-omp-ext          \

// DEFINE:          -ferror-limit=100 %t-acc.c |                               \

// DEFINE:   FileCheck -check-prefixes=%{PRT_CHK} -match-full-lines            \

// DEFINE:             -strict-whitespace %s

/* noerrs-no-diagnostics */

//      PRT:int i;

// PRT-NEXT:int non_const_expr = 2;

// PRT-NEXT:int possibleSideEffects() {

// PRT-NEXT:  return 2;

// PRT-NEXT:}

int i;

int non_const_expr = 2;

int possibleSideEffects() {

  return 2;

}

//--------------------------------------------------

//------------------------------------------------------------------------------

// Translatable directives are translated before any error.

//--------------------------------------------------

//------------------------------------------------------------------------------

// PRT-NEXT:void beforeError() {

void beforeError() {

#pragma acc routine seq

void beforeError_routine();

//--------------------------------------------------

//------------------------------------------------------------------------------

// Directive in macro expansion: yes

// No associated code.

//--------------------------------------------------

//------------------------------------------------------------------------------

// PRT-NEXT:void inMacroNoAssoc() {

void inMacroNoAssoc() {

  #undef MAC

}// PRT-NEXT:}

//--------------------------------------------------

//------------------------------------------------------------------------------

// Directive in macro expansion: yes

// Associated code ends in macro expansion: no

//--------------------------------------------------

//------------------------------------------------------------------------------

// PRT-NEXT:void inMacroAssocEndNotInMacro() {

void inMacroAssocEndNotInMacro() {

void inMacroAssocEndNotInMacro_routine();

#undef MAC

//--------------------------------------------------

//------------------------------------------------------------------------------

// Directive in macro expansion: yes

// Associated code ends in macro expansion: yes

//

// exercise different code paths, so check all of them.  For expression

// statements, whether the token preceding the semicolon is part of the same

// macro expansion also affects the code path.

//--------------------------------------------------

//------------------------------------------------------------------------------

// PRT-NEXT:void inMacroAssocEndInMacro() {

void inMacroAssocEndInMacro() {

#undef MAC2

#undef MAC3

//--------------------------------------------------

//------------------------------------------------------------------------------

// Directive in macro expansion: no

// No associated code.

//--------------------------------------------------

//------------------------------------------------------------------------------

// PRT-NEXT:void notInMacroNoAssoc() {

void notInMacroNoAssoc() {

  _Pragma("acc update device(i)")

}// PRT-NEXT:}

//--------------------------------------------------

//------------------------------------------------------------------------------

// Directive in macro expansion: no

// Associated code ends in macro expansion: no

//--------------------------------------------------

//------------------------------------------------------------------------------

// PRT-NEXT:/* expected-error{{.*}} */

// PRT-NEXT:_Pragma("acc routine seq")

_Pragma("acc routine seq")

void notInMacroAssocEndNotInMacro_routine();

//--------------------------------------------------

//------------------------------------------------------------------------------

// Directive in macro expansion: no

// Associated code ends in macro expansion: yes

//

// so check all of them.  For expression statements and function prototypes,

// whether the token preceding the semicolon is part of the same macro expansion

// also affects the code path.

//--------------------------------------------------

//------------------------------------------------------------------------------

// PRT-NEXT:void notInMacroAssocEndInMacro() {

void notInMacroAssocEndInMacro() {

  // PRT-NEXT:  #define MAC ;

  // PRT-NEXT:  #pragma acc parallel num_workers(non_const_expr)

  // PRT-NEXT:  #pragma acc loop worker

  // PRT-NEXT:  #pragma acc parallel vector_length(possibleSideEffects())

  // PRT-NEXT:  #pragma acc loop vector

  // PRT-NEXT:  for (int i = 0; i < 5; ++i)

  // PRT-NEXT:    /* expected-error{{.*}} */

  // PRT-NEXT:    MAC

  // PRT-NEXT:  #undef MAC

  #define MAC ;

  #pragma acc parallel num_workers(non_const_expr)

  #pragma acc loop worker

  #pragma acc parallel vector_length(possibleSideEffects())

  #pragma acc loop vector

  for (int i = 0; i < 5; ++i)

    /* expected-error@+1 {{cannot rewrite OpenACC construct ending within a macro expansion}} */

    MAC

  #undef MAC

  // PRT-NEXT:  #define MAC {}

  // PRT-NEXT:  #pragma acc parallel num_workers(non_const_expr)

  // PRT-NEXT:  #pragma acc loop worker

  // PRT-NEXT:  #pragma acc parallel vector_length(possibleSideEffects())

  // PRT-NEXT:  #pragma acc loop vector

  // PRT-NEXT:  for (int i = 0; i < 5; ++i)

  // PRT-NEXT:    /* expected-error{{.*}} */

  // PRT-NEXT:    MAC

  // PRT-NEXT:  #undef MAC

  #define MAC {}

  #pragma acc parallel num_workers(non_const_expr)

  #pragma acc loop worker

  #pragma acc parallel vector_length(possibleSideEffects())

  #pragma acc loop vector

  for (int i = 0; i < 5; ++i)

    /* expected-error@+1 {{cannot rewrite OpenACC construct ending within a macro expansion}} */

    MAC

  #undef MAC

  // PRT-NEXT:  #define MAC }

  // PRT-NEXT:  #pragma acc parallel num_workers(non_const_expr)

  // PRT-NEXT:  #pragma acc loop worker

  // PRT-NEXT:  #pragma acc parallel vector_length(possibleSideEffects())

  // PRT-NEXT:  #pragma acc loop vector

  // PRT-NEXT:  for (int i = 0; i < 5; ++i)

  // PRT-NEXT:    /* expected-error{{.*}} */

  // PRT-NEXT:    {MAC

  // PRT-NEXT:  #undef MAC

  #define MAC }

  #pragma acc parallel num_workers(non_const_expr)

  #pragma acc loop worker

  #pragma acc parallel vector_length(possibleSideEffects())

  #pragma acc loop vector

  for (int i = 0; i < 5; ++i)

    /* expected-error@+1 {{cannot rewrite OpenACC construct ending within a macro expansion}} */

    {MAC

  #undef MAC

  // PRT-NEXT:  #define MAC (i = 3);

  // PRT-NEXT:  #pragma acc parallel num_workers(non_const_expr)

  // PRT-NEXT:  #pragma acc loop worker

  // PRT-NEXT:  #pragma acc parallel vector_length(possibleSideEffects())

  // PRT-NEXT:  #pragma acc loop vector

  // PRT-NEXT:  for (int j = 0; j < 5; ++j)

  // PRT-NEXT:    /* expected-error{{.*}} */

  // PRT-NEXT:    MAC

  // PRT-NEXT:  #undef MAC

  #define MAC (i = 3);

  #pragma acc parallel num_workers(non_const_expr)

  #pragma acc loop worker

  #pragma acc parallel vector_length(possibleSideEffects())

  #pragma acc loop vector

  for (int j = 0; j < 5; ++j)

    /* expected-error@+1 {{cannot rewrite OpenACC construct ending within a macro expansion}} */

    MAC

  #undef MAC

  // PRT-NEXT:  #define MAC ;

  // PRT-NEXT:  #pragma acc parallel num_workers(non_const_expr)

  // PRT-NEXT:  #pragma acc loop worker

  // PRT-NEXT:  #pragma acc parallel vector_length(possibleSideEffects())

  // PRT-NEXT:  #pragma acc loop vector

  // PRT-NEXT:  for (int j = 0; j < 5; ++j)

  // PRT-NEXT:    /* expected-error{{.*}} */

  // PRT-NEXT:    (i = 3) MAC

  // PRT-NEXT:  #undef MAC

  #define MAC ;

  #pragma acc parallel num_workers(non_const_expr)

  #pragma acc loop worker

  #pragma acc parallel vector_length(possibleSideEffects())

  #pragma acc loop vector

  for (int j = 0; j < 5; ++j)

    /* expected-error@+1 {{cannot rewrite OpenACC construct ending within a macro expansion}} */

    (i = 3) MAC

  // PRT-NEXT:  #define MAC1 )

  // PRT-NEXT:  #define MAC2 ;

  // PRT-NEXT:  #pragma acc parallel num_workers(non_const_expr)

  // PRT-NEXT:  #pragma acc loop worker

  // PRT-NEXT:  #pragma acc parallel vector_length(possibleSideEffects())

  // PRT-NEXT:  #pragma acc loop vector

  // PRT-NEXT:  for (int j = 0; j < 5; ++j)

  // PRT-NEXT:    /* expected-error{{.*}} */

  // PRT-NEXT:    (i = 3 MAC1 MAC2

  // PRT-NEXT:  #undef MAC2

  #define MAC1 )

  #define MAC2 ;

  #pragma acc parallel num_workers(non_const_expr)

  #pragma acc loop worker

  #pragma acc parallel vector_length(possibleSideEffects())

  #pragma acc loop vector

  for (int j = 0; j < 5; ++j)

    /* expected-error@+1 {{cannot rewrite OpenACC construct ending within a macro expansion}} */

    (i = 3 MAC1 MAC2

void notInMacroAssocEndInMacro_routineProtoLastToken() MAC

#undef MAC

//--------------------------------------------------

//------------------------------------------------------------------------------

// Translatable directives are translated after any error.

//--------------------------------------------------

//------------------------------------------------------------------------------

// PRT-NEXT:void afterError() {

void afterError() {