[flang][OpenMP] Add semantic check for target nesting

[peixin]

Summary

This patch implements the following check for TARGET construct:

OpenMP Version 5.0 Target construct restriction: If a target update,
target data, target enter data, or target exit data construct is
encountered during execution of a target region, the behavior is
unspecified.

Also add one test case for the check.

[clementval]

flang/lib/Semantics/check-omp-structure.cpp

291

It would be better to have a generic IsNestedIn function where you can pass the directive you want to check for. We don't want to have a counter for each construct.

516

The standard says it is unspecified but not really restricted.

668

As mentioned above you should use the current _dirContext stack with a generic function.

738

Same here. Please use generic mechanism.

flang/test/Semantics/omp-target01.f90

1 ↗

(On Diff #359375)

You don't check if your restriction actually work here.

[peixin]

@clementval Thanks for the comments. Upload the following update:

1. I am reminded that the counter is not necessary since the if branch dirContext_.size() >= 1 already checks the directive nested. Delete the counter usage and checking GetContext().directive == llvm::omp::Directive::OMPD_target inside dirContext_.size() >= 1 check is enough.
2. Add the error check in test case.
3. There is no this error info before, or at lease I did not find in f18 openmp semantic check. I test gfortran 9.3.0 and it gives the following warning:

   15 |   !$omp target update from(arrayA) to(arrayB)
Warning: ‘target update’ construct inside of ‘target’ region
   23 |   !$omp target data map(to: a)
Warning: ‘target data’ construct inside of ‘target’ region
   33 |   !$omp target enter data map(alloc:B)
Warning: ‘target enter data’ construct inside of ‘target’ region
   38 |   !$omp target exit data map(delete:B)
Warning: ‘target exit data’ construct inside of ‘target’ region

I also generate one similar c code to check target data directive according to OpenMP API Examples Version 5.0.1 as follows:

extern void init(float*, float*, int);
extern void output(float*, int);
void vec_mult(float *p, float *v1, float *v2, int N)
{
  int i;
  init(v1, v2, N);
  #pragma omp target
  #pragma omp target data map(to: v1[0:N], v2[:N]) map(from: p[0:N])
  {
    for (i=0; i<N; i++)
      p[i] = v1[i] * v2[i];
  }
  output(p, N);
}

Tested using gcc 9.3.0 and current clang and the results are as follows:

$ clang -fopenmp temp.c 
temp.c:8:3: error: region cannot be nested inside 'target' region
  #pragma omp target data map(to: v1[0:N], v2[:N]) map(from: p[0:N])
  ^
1 error generated.
$ gcc -fopenmp temp.c 
temp.c: In function ‘vec_mult’:
temp.c:8:11: warning: ‘target data’ construct inside of ‘target’ region
    8 |   #pragma omp target data map(to: v1[0:N], v2[:N]) map(from: p[0:N])
      |           ^~~

Then I use the similar error info as clang to keep the consistence.

[peixin]

Set the flag eligibleTARGET as true by default and print the error when it is false.

[kiranchandramohan]

flang/lib/Semantics/check-omp-structure.cpp

479

Is the restriction for closely nested or does it apply to any kind of nesting? For e.g is the following allowed and will the test catch this?

-> target
-> -> parallel
-> -> -> target data

516

Same concern as @clementval. Should this be a warning or an error?

flang/lib/Semantics/check-omp-structure.h

228

Nit: I think it will be better to rename this as CheckTargetNest.

[peixin]

@kiranchandramohan Thanks for the review.

The restriction applies to any kind of nesting. Add the test scenario of parallel directive between target and target update, which is not covered in previous design. Fixed the design to cover this scenario. Also rename the function name.

The error reporting is truly not completely reasonable. Change it into warning report and also fix the error info by replacing cannot by should not. Is this OK with you? @kiranchandramohan @clementval

[peixin]

flang/lib/Semantics/check-omp-structure.cpp

291

I am thinking about if we can replace this for loop search with the counter marked for target nesting since the for loop is executed for each openmp construct check, which is not efficient. With the design in https://reviews.llvm.org/D106335, I only need add one more element in the enum definition of directiveNestType. What do you think?

[kiranchandramohan]

Have some suggestions.

flang/lib/Semantics/check-omp-structure.cpp

291

That sounds good to me.

I think ideally, you can do the following.
On visiting a target_data, target_update, target_enter_data or target_exit_data directive check if it is nested in a target region using GetDirectiveNest and report a warning if it is nested.

481

The name of the variable does not match its intention. Like these are the ones that cause a warning if true, right?

519

Can the warning be changed to what the standard says, that the "behaviour is unspecified"?

[peixin]

I will rebase this patch or https://reviews.llvm.org/D106335 after either of them is merged.

[peixin]

@kiranchandramohan Thanks for the comments.

flang/lib/Semantics/check-omp-structure.cpp

291

Thanks for the suggestion. Fixed now.

481

Yes. Sorry for that. Reversed the value to match its intention.

519

Replaced the warning message with what the standard says according to your suggestion.

[kiranchandramohan]

Looks OK. I have a question and a suggestion. Please make the change if you agree otherwise let me know.

flang/lib/Semantics/check-omp-structure.cpp

291

I was suggesting something like the following to avoid double traversal of OpenMPBlockConstruct. We would be visiting these nodes anyway during regular traversal.
Have you considered this? Feel free to dissent.

void OmpStructureChecker::Enter(const parser::OpenMPBlockConstruct &x) {
  const auto &beginBlockDir{std::get<parser::OmpBeginBlockDirective>(c.t)};
  const auto &beginDir{std::get<parser::OmpBlockDirective>(beginBlockDir.t)};
   if (GetDirectiveNest(TargetNest) > 0) {
     if (beginDir.v == llvm::omp::Directive::OMPD_target_data) {
       //print warning
     }
}

481

Nit:Would eligibleTargetDir be better?

[peixin]

@kiranchandramohan Thanks for the review.

flang/lib/Semantics/check-omp-structure.cpp

291

Actually, my first implementation is like what you suggest. Arnamoy suggests the current implementation and I argee with him. Merging the semantic checks for target_data and other three directives in one function seems easier to read and understand. In addition, implementing the semantic check of the four directives nested inside target in one fuction seems more reasonable since they are in one restriction. This is one time-readability tradeoff. I think the double traversal does not take too much time. What do you think?

481

elibleTargetDir does not match its intention. Would ineligibleTargetDir be better?

[kiranchandramohan]

LGTM. One minor comment.

@clementval do you have further comments?

flang/lib/Semantics/check-omp-structure.cpp

291

OK.

481

Yes, I would prefer ineligibleTargetDir.

[peixin]

@kiranchandramohan Thanks for the review.
Replaced the variable name directive with ineligibleTargetDir.

[peixin]

Fix the clang-format warning.

[clementval]

LGTM. Sorry I was kinda busy the last few days and forgot to review the changes. Please go ahead and land the patch.

[peixin]

Rebase this patch.

[peixin]

Fix the clang format warning.

flang/lib/Semantics/check-directive-structure.h

protected:
   void SetContextClauseInfo(C type) {
     GetContext().clauseInfo.emplace(type, GetContext().clause);
   }
 
   void AddClauseToCrtContext(C type) {
     GetContext().actualClauses.push_back(type);
   }
 
-  void EnterSIMDNest() { simdNest_++; }
-
-  void ExitSIMDNest() { simdNest_--; }
-
-  int GetSIMDNest() { return simdNest_; }
-
   // Check if the given clause is present in the current context
   const PC *FindClause(C type) {
     auto it{GetContext().clauseInfo.find(type)};
     if (it != GetContext().clauseInfo.end()) {
       return it->second;
     }
     return nullptr;
   }
protected:
   virtual llvm::StringRef getDirectiveName(D directive) { return ""; };
 
   SemanticsContext &context_;
   std::vector<DirectiveContext> dirContext_; // used as a stack
   std::unordered_map<D, DirectiveClauses<C, ClauseEnumSize>>
       directiveClausesMap_;
 
   std::string ClauseSetToString(const common::EnumSet<C, ClauseEnumSize> set);
-  int simdNest_{0};
 };
 
 template <typename D, typename C, typename PC, std::size_t ClauseEnumSize>
 void DirectiveStructureChecker<D, C, PC, ClauseEnumSize>::CheckNoBranching(
     const parser::Block &block, D directive,
     const parser::CharBlock &directiveSource) {
   NoBranchingEnforce<D> noBranchingEnforce{
       context_, directiveSource, directive, ContextDirectiveAsFortran()};

flang/lib/Semantics/check-omp-structure.h

void CheckIsLoopIvPartOfClause(
       llvmOmpClause clause, const parser::OmpObjectList &ompObjectList);
   void CheckWorkshareBlockStmts(const parser::Block &, parser::CharBlock);
 
   void CheckLoopItrVariableIsInt(const parser::OpenMPLoopConstruct &x);
   void CheckDoWhile(const parser::OpenMPLoopConstruct &x);
   void CheckCycleConstraints(const parser::OpenMPLoopConstruct &x);
   void CheckDistLinear(const parser::OpenMPLoopConstruct &x);
   void CheckSIMDNest(const parser::OpenMPConstruct &x);
+  void CheckTargetNest(const parser::OpenMPConstruct &x);
   std::int64_t GetOrdCollapseLevel(const parser::OpenMPLoopConstruct &x);
   void CheckIfDoOrderedClause(const parser::OmpBlockDirective &blkDirectiv);
   bool CheckReductionOperators(const parser::OmpClause::Reduction &);
   bool CheckIntrinsicOperator(
       const parser::DefinedOperator::IntrinsicOperator &);
   void CheckReductionTypeList(const parser::OmpClause::Reduction &);
   void CheckMasterNesting(const parser::OpenMPBlockConstruct &x);
   void CheckBarrierNesting(const parser::OpenMPSimpleStandaloneConstruct &x);
   void CheckReductionArraySection(const parser::OmpObjectList &ompObjectList);
   void CheckIntentInPointerAndDefinable(
       const parser::OmpObjectList &, const llvm::omp::Clause);
   void CheckArraySection(const parser::ArrayElement &arrayElement,
       const parser::Name &name, const llvm::omp::Clause clause);
   void CheckMultipleAppearanceAcrossContext(
       const parser::OmpObjectList &ompObjectList);
   const parser::OmpObjectList *GetOmpObjectList(const parser::OmpClause &);
   void CheckPredefinedAllocatorRestriction(const parser::CharBlock &source,
       const parser::OmpObjectList &ompObjectList);
   void CheckPredefinedAllocatorRestriction(
       const parser::CharBlock &source, const parser::Name &name);
   bool isPredefinedAllocator{false};
+  void EnterDirectiveNest(const int index) { directiveNest_[index]++; }
+  void ExitDirectiveNest(const int index) { directiveNest_[index]--; }
+  int GetDirectiveNest(const int index) { return directiveNest_[index]; }
+
+  enum directiveNestType { SIMDNest, TargetNest, LastType };
+  int directiveNest_[LastType + 1] = {0};
 };
 } // namespace Fortran::semantics
 #endif // FORTRAN_SEMANTICS_CHECK_OMP_STRUCTURE_H_

flang/lib/Semantics/check-omp-structure.cpp

 
 void OmpStructureChecker::Enter(const parser::OpenMPConstruct &x) {
   // Simd Construct with Ordered Construct Nesting check
   // We cannot use CurrentDirectiveIsNested() here because
   // PushContextAndClauseSets() has not been called yet, it is
   // called individually for each construct.  Therefore a
   // dirContext_ size `1` means the current construct is nested
   if (dirContext_.size() >= 1) {
-    if (GetSIMDNest() > 0) {
+    if (GetDirectiveNest(SIMDNest) > 0) {
       CheckSIMDNest(x);
     }
+    if (GetDirectiveNest(TargetNest) > 0) {
+      CheckTargetNest(x);
+    }
   }
 }
 
 void OmpStructureChecker::Enter(const parser::OpenMPLoopConstruct &x) {
   const auto &beginLoopDir{std::get<parser::OmpBeginLoopDirective>(x.t)};
   const auto &beginDir{std::get<parser::OmpLoopDirective>(beginLoopDir.t)};
 
   // check matching, End directive is optional
   if (const auto &endLoopDir{
           std::get<std::optional<parser::OmpEndLoopDirective>>(x.t)}) {
     const auto &endDir{
         std::get<parser::OmpLoopDirective>(endLoopDir.value().t)};
 
     CheckMatching<parser::OmpLoopDirective>(beginDir, endDir);
   }
 
   PushContextAndClauseSets(beginDir.source, beginDir.v);
   if (llvm::omp::simdSet.test(GetContext().directive)) {
-    EnterSIMDNest();
+    EnterDirectiveNest(SIMDNest);
   }
 
   if (beginDir.v == llvm::omp::Directive::OMPD_do) {
     // 2.7.1 do-clause -> private-clause |
     //                    firstprivate-clause |
     //                    lastprivate-clause |
     //                    linear-clause |
     //                    reduction-clause |
if (!eligibleSIMD) {
     context_.Say(parser::FindSourceLocation(c),
         "The only OpenMP constructs that can be encountered during execution "
         "of a 'SIMD'"
         " region are the `ATOMIC` construct, the `LOOP` construct, the `SIMD`"
         " construct and the `ORDERED` construct with the `SIMD` clause."_err_en_US);
   }
 }
 
+void OmpStructureChecker::CheckTargetNest(const parser::OpenMPConstruct &c) {
+  // 2.12.5 Target Construct Restriction
+  bool eligibleTarget{true};
+  llvm::omp::Directive directive;
+  std::visit(
+      common::visitors{
+          [&](const parser::OpenMPBlockConstruct &c) {
+            const auto &beginBlockDir{
+                std::get<parser::OmpBeginBlockDirective>(c.t)};
+            const auto &beginDir{
+                std::get<parser::OmpBlockDirective>(beginBlockDir.t)};
+            if (beginDir.v == llvm::omp::Directive::OMPD_target_data) {
+              eligibleTarget = false;
+              directive = beginDir.v;
+            }
+          },
+          [&](const parser::OpenMPStandaloneConstruct &c) {
+            std::visit(
+                common::visitors{
+                    [&](const parser::OpenMPSimpleStandaloneConstruct &c) {
+                      const auto &dir{
+                          std::get<parser::OmpSimpleStandaloneDirective>(c.t)};
+                      if (dir.v == llvm::omp::Directive::OMPD_target_update ||
+                          dir.v ==
+                              llvm::omp::Directive::OMPD_target_enter_data ||
+                          dir.v ==
+                              llvm::omp::Directive::OMPD_target_exit_data) {
+                        eligibleTarget = false;
+                        directive = dir.v;
+                      }
+                    },
+                    [&](const auto &c) {},
+                },
+                c.u);
+          },
+          [&](const auto &c) {},
+      },
+      c.u);
+  if (!eligibleTarget) {
+    context_.Say(parser::FindSourceLocation(c),
+        "If %s directive is nested inside TARGET region, the behaviour "
+        "is unspecified"_en_US,
+        parser::ToUpperCaseLetters(getDirectiveName(directive).str()));
+  }
+}
+
 std::int64_t OmpStructureChecker::GetOrdCollapseLevel(
     const parser::OpenMPLoopConstruct &x) {
   const auto &beginLoopDir{std::get<parser::OmpBeginLoopDirective>(x.t)};
   const auto &clauseList{std::get<parser::OmpClauseList>(beginLoopDir.t)};
   std::int64_t orderedCollapseLevel{1};
   std::int64_t orderedLevel{0};
   std::int64_t collapseLevel{0};
 
for (auto var : indexVars) {
           "Variable '%s' not allowed in `LINEAR` clause, only loop iterator can be specified in `LINEAR` clause of a construct combined with `DISTRIBUTE`"_err_en_US,
           root.name());
     }
   }
 }
 
 void OmpStructureChecker::Leave(const parser::OpenMPLoopConstruct &) {
   if (llvm::omp::simdSet.test(GetContext().directive)) {
-    ExitSIMDNest();
+    ExitDirectiveNest(SIMDNest);
   }
   dirContext_.pop_back();
 }
 
 void OmpStructureChecker::Enter(const parser::OmpEndLoopDirective &x) {
   const auto &dir{std::get<parser::OmpLoopDirective>(x.t)};
   ResetPartialContext(dir.source);
   switch (dir.v) {
void OmpStructureChecker::Enter(const parser::OpenMPBlockConstruct &x) {
   const auto &endBlockDir{std::get<parser::OmpEndBlockDirective>(x.t)};
   const auto &beginDir{std::get<parser::OmpBlockDirective>(beginBlockDir.t)};
   const auto &endDir{std::get<parser::OmpBlockDirective>(endBlockDir.t)};
   const parser::Block &block{std::get<parser::Block>(x.t)};
 
   CheckMatching<parser::OmpBlockDirective>(beginDir, endDir);
 
   PushContextAndClauseSets(beginDir.source, beginDir.v);
+  if (GetContext().directive == llvm::omp::Directive::OMPD_target) {
+    EnterDirectiveNest(TargetNest);
+  }
 
   if (CurrentDirectiveIsNested()) {
     CheckIfDoOrderedClause(beginDir);
     if (llvm::omp::teamSet.test(GetContextParent().directive)) {
       HasInvalidTeamsNesting(beginDir.v, beginDir.source);
     }
     if (GetContext().directive == llvm::omp::Directive::OMPD_master) {
       CheckMasterNesting(x);
else if (IsCloselyNestedRegion(llvm::omp::nestedOrderedErrSet)) {
       context_.Say(blkDirective.source,
           "`ORDERED` region may not be closely nested inside of "
           "`CRITICAL`, `ORDERED`, explicit `TASK` or `TASKLOOP` region."_err_en_US);
     }
   }
 }
 
 void OmpStructureChecker::Leave(const parser::OpenMPBlockConstruct &) {
+  if (GetContext().directive == llvm::omp::Directive::OMPD_target) {
+    ExitDirectiveNest(TargetNest);
+  }
   dirContext_.pop_back();
 }
 
 void OmpStructureChecker::Enter(const parser::OpenMPSectionsConstruct &x) {
   const auto &beginSectionsDir{
       std::get<parser::OmpBeginSectionsDirective>(x.t)};
   const auto &endSectionsDir{std::get<parser::OmpEndSectionsDirective>(x.t)};
   const auto &beginDir{

flang/test/Semantics/omp-nested-target.f90

+! RUN: %S/test_errors.sh %s %t %flang_fc1 -fopenmp
+! REQUIRES: shell
+
+! OpenMP Version 5.0
+! Check OpenMP construct validity for the following directives:
+! 2.12.5 Target Construct
+
+program main
+  integer :: i, j, N = 10
+  real :: a, arrayA(512), arrayB(512), ai(10)
+  real, allocatable :: B(:)
+
+  !$omp target
+  !WARNING: If TARGET UPDATE directive is nested inside TARGET region, the behaviour is unspecified
+  !$omp target update from(arrayA) to(arrayB)
+  do i = 1, 512
+    arrayA(i) = arrayB(i)
+  end do
+  !$omp end target
+
+  !$omp parallel
+  !$omp target
+  !$omp parallel
+  !WARNING: If TARGET UPDATE directive is nested inside TARGET region, the behaviour is unspecified
+  !$omp target update from(arrayA) to(arrayB)
+  do i = 1, 512
+    arrayA(i) = arrayB(i)
+  end do
+  !$omp end parallel
+  !$omp end target
+  !$omp end parallel
+
+  !$omp target
+  !WARNING: If TARGET DATA directive is nested inside TARGET region, the behaviour is unspecified
+  !$omp target data map(to: a)
+  do i = 1, N
+    a = 3.14
+  end do
+  !$omp end target data
+  !$omp end target
+
+  allocate(B(N))
+  !$omp target
+  !WARNING: If TARGET ENTER DATA directive is nested inside TARGET region, the behaviour is unspecified
+  !$omp target enter data map(alloc:B)
+  !$omp end target
+
+  !$omp target
+  !WARNING: If TARGET EXIT DATA directive is nested inside TARGET region, the behaviour is unspecified
+  !$omp target exit data map(delete:B)
+  !$omp end target
+  deallocate(B)
+
+end program main