Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[flang][OpenMP] Add semantic check for declare target #72770

Merged
merged 1 commit into from
Nov 22, 2023

Conversation

shraiysh
Copy link
Member

This patch adds the following check from OpenMP 5.2.

If the directive has a clause, it must contain at least one enter clause
or at least one link clause.

Also added a warning for the deprication of TO clause on DECLARE TARGET construct.

The clause-name to may be used as a synonym for the clause-name enter.
This use has been deprecated.

Last attempt to add this caused issues in buildbots - #71861. It should pass now.

@llvmbot llvmbot added clang Clang issues not falling into any other category clang:frontend Language frontend issues, e.g. anything involving "Sema" flang Flang issues not falling into any other category flang:fir-hlfir flang:openmp flang:semantics flang:parser clang:openmp OpenMP related changes to Clang labels Nov 19, 2023
@llvmbot
Copy link
Collaborator

llvmbot commented Nov 19, 2023

@llvm/pr-subscribers-clang
@llvm/pr-subscribers-flang-semantics
@llvm/pr-subscribers-flang-openmp
@llvm/pr-subscribers-flang-parser

@llvm/pr-subscribers-flang-fir-hlfir

Author: Shraiysh (shraiysh)

Changes

This patch adds the following check from OpenMP 5.2.

If the directive has a clause, it must contain at least one enter clause
or at least one link clause.

Also added a warning for the deprication of TO clause on DECLARE TARGET construct.

The clause-name to may be used as a synonym for the clause-name enter.
This use has been deprecated.

Last attempt to add this caused issues in buildbots - #71861. It should pass now.


Patch is 67.44 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/72770.diff

24 Files Affected:

  • (modified) clang/lib/Parse/ParseOpenMP.cpp (+1)
  • (modified) clang/test/OpenMP/target_enter_data_nowait_messages.cpp (+13-1)
  • (modified) flang/lib/Lower/OpenMP.cpp (+15)
  • (modified) flang/lib/Parser/openmp-parsers.cpp (+2)
  • (modified) flang/lib/Semantics/check-omp-structure.cpp (+39-12)
  • (modified) flang/lib/Semantics/check-omp-structure.h (+2)
  • (modified) flang/lib/Semantics/resolve-directives.cpp (+3)
  • (modified) flang/test/Lower/OpenMP/FIR/declare-target-data.f90 (+16)
  • (modified) flang/test/Lower/OpenMP/FIR/declare-target-func-and-subr.f90 (+68)
  • (added) flang/test/Lower/OpenMP/FIR/declare-target-implicit-func-and-subr-cap-enter.f90 (+192)
  • (modified) flang/test/Lower/OpenMP/declare-target-data.f90 (+16)
  • (modified) flang/test/Lower/OpenMP/declare-target-func-and-subr.f90 (+68)
  • (added) flang/test/Lower/OpenMP/declare-target-implicit-func-and-subr-cap-enter.f90 (+192)
  • (modified) flang/test/Lower/OpenMP/declare-target-implicit-tarop-cap.f90 (+14)
  • (modified) flang/test/Lower/OpenMP/function-filtering-2.f90 (+8)
  • (modified) flang/test/Lower/OpenMP/function-filtering.f90 (+18)
  • (modified) flang/test/Parser/OpenMP/declare_target-device_type.f90 (+21-6)
  • (modified) flang/test/Semantics/OpenMP/declarative-directive.f90 (+11-1)
  • (modified) flang/test/Semantics/OpenMP/declare-target01.f90 (+48)
  • (modified) flang/test/Semantics/OpenMP/declare-target02.f90 (+54)
  • (modified) flang/test/Semantics/OpenMP/declare-target06.f90 (+5)
  • (modified) flang/test/Semantics/OpenMP/requires04.f90 (+4-1)
  • (modified) flang/test/Semantics/OpenMP/requires05.f90 (+2)
  • (modified) llvm/include/llvm/Frontend/OpenMP/OMP.td (+5-1)
diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp
index 3e7d8274aeefc52..ff3417f5a251ade 100644
--- a/clang/lib/Parse/ParseOpenMP.cpp
+++ b/clang/lib/Parse/ParseOpenMP.cpp
@@ -3369,6 +3369,7 @@ OMPClause *Parser::ParseOpenMPClause(OpenMPDirectiveKind DKind,
   case OMPC_exclusive:
   case OMPC_affinity:
   case OMPC_doacross:
+  case OMPC_enter:
     if (getLangOpts().OpenMP >= 52 && DKind == OMPD_ordered &&
         CKind == OMPC_depend)
       Diag(Tok, diag::warn_omp_depend_in_ordered_deprecated);
diff --git a/clang/test/OpenMP/target_enter_data_nowait_messages.cpp b/clang/test/OpenMP/target_enter_data_nowait_messages.cpp
index 3f5dde00b814ca4..ba5eaf1d2214a07 100644
--- a/clang/test/OpenMP/target_enter_data_nowait_messages.cpp
+++ b/clang/test/OpenMP/target_enter_data_nowait_messages.cpp
@@ -6,14 +6,26 @@ int main(int argc, char **argv) {
   int i;
 
   #pragma omp nowait target enter data map(to: i) // expected-error {{expected an OpenMP directive}}
-  #pragma omp target nowait enter data map(to: i) // expected-warning {{extra tokens at the end of '#pragma omp target' are ignored}}
+  {}
+  #pragma omp target nowait foo data map(to: i) // expected-warning {{extra tokens at the end of '#pragma omp target' are ignored}}
+  {}
+  #pragma omp target nowait enter data map(to: i) // expected-warning {{extra tokens at the end of '#pragma omp target' are ignored}} expected-error {{unexpected OpenMP clause 'enter' in directive '#pragma omp target'}} expected-error {{expected '(' after 'enter'}}
+  {}
   #pragma omp target enter nowait data map(to: i) // expected-error {{expected an OpenMP directive}}
+  {}
   #pragma omp target enter data nowait() map(to: i) // expected-warning {{extra tokens at the end of '#pragma omp target enter data' are ignored}} expected-error {{expected at least one 'map' clause for '#pragma omp target enter data'}}
+  {}
   #pragma omp target enter data map(to: i) nowait( // expected-warning {{extra tokens at the end of '#pragma omp target enter data' are ignored}}
+  {}
   #pragma omp target enter data map(to: i) nowait (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target enter data' are ignored}}
+  {}
   #pragma omp target enter data map(to: i) nowait device (-10u)
+  {}
   #pragma omp target enter data map(to: i) nowait (3.14) device (-10u) // expected-warning {{extra tokens at the end of '#pragma omp target enter data' are ignored}}
+  {}
   #pragma omp target enter data map(to: i) nowait nowait // expected-error {{directive '#pragma omp target enter data' cannot contain more than one 'nowait' clause}}
+  {}
   #pragma omp target enter data nowait map(to: i) nowait // expected-error {{directive '#pragma omp target enter data' cannot contain more than one 'nowait' clause}}
+  {}
   return 0;
 }
diff --git a/flang/lib/Lower/OpenMP.cpp b/flang/lib/Lower/OpenMP.cpp
index 00f16026d9b4bb8..6657f46ba6325ca 100644
--- a/flang/lib/Lower/OpenMP.cpp
+++ b/flang/lib/Lower/OpenMP.cpp
@@ -590,6 +590,8 @@ class ClauseProcessor {
   bool processSectionsReduction(mlir::Location currentLocation) const;
   bool processTo(llvm::SmallVectorImpl<DeclareTargetCapturePair> &result) const;
   bool
+  processEnter(llvm::SmallVectorImpl<DeclareTargetCapturePair> &result) const;
+  bool
   processUseDeviceAddr(llvm::SmallVectorImpl<mlir::Value> &operands,
                        llvm::SmallVectorImpl<mlir::Type> &useDeviceTypes,
                        llvm::SmallVectorImpl<mlir::Location> &useDeviceLocs,
@@ -1851,6 +1853,18 @@ bool ClauseProcessor::processTo(
       });
 }
 
+bool ClauseProcessor::processEnter(
+    llvm::SmallVectorImpl<DeclareTargetCapturePair> &result) const {
+  return findRepeatableClause<ClauseTy::Enter>(
+      [&](const ClauseTy::Enter *enterClause,
+          const Fortran::parser::CharBlock &) {
+        // Case: declare target to(func, var1, var2)...
+        gatherFuncAndVarSyms(enterClause->v,
+                             mlir::omp::DeclareTargetCaptureClause::enter,
+                             result);
+      });
+}
+
 bool ClauseProcessor::processUseDeviceAddr(
     llvm::SmallVectorImpl<mlir::Value> &operands,
     llvm::SmallVectorImpl<mlir::Type> &useDeviceTypes,
@@ -2792,6 +2806,7 @@ static mlir::omp::DeclareTargetDeviceType getDeclareTargetInfo(
 
     ClauseProcessor cp(converter, *clauseList);
     cp.processTo(symbolAndClause);
+    cp.processEnter(symbolAndClause);
     cp.processLink(symbolAndClause);
     cp.processDeviceType(deviceType);
     cp.processTODO<Fortran::parser::OmpClause::Indirect>(
diff --git a/flang/lib/Parser/openmp-parsers.cpp b/flang/lib/Parser/openmp-parsers.cpp
index 0271f699a687ff0..bba1be27158ce7e 100644
--- a/flang/lib/Parser/openmp-parsers.cpp
+++ b/flang/lib/Parser/openmp-parsers.cpp
@@ -259,6 +259,8 @@ TYPE_PARSER(
             parenthesized("STATIC" >> maybe("," >> scalarIntExpr)))) ||
     "DYNAMIC_ALLOCATORS" >>
         construct<OmpClause>(construct<OmpClause::DynamicAllocators>()) ||
+    "ENTER" >> construct<OmpClause>(construct<OmpClause::Enter>(
+                   parenthesized(Parser<OmpObjectList>{}))) ||
     "FINAL" >> construct<OmpClause>(construct<OmpClause::Final>(
                    parenthesized(scalarLogicalExpr))) ||
     "FULL" >> construct<OmpClause>(construct<OmpClause::Full>()) ||
diff --git a/flang/lib/Semantics/check-omp-structure.cpp b/flang/lib/Semantics/check-omp-structure.cpp
index 0b1a581bf298196..4dbbf1260f3ea31 100644
--- a/flang/lib/Semantics/check-omp-structure.cpp
+++ b/flang/lib/Semantics/check-omp-structure.cpp
@@ -1165,13 +1165,31 @@ void OmpStructureChecker::Enter(const parser::OmpClause::Allocate &x) {
   }
 }
 
+void OmpStructureChecker::Enter(const parser::OmpDeclareTargetWithClause &x) {
+  SetClauseSets(llvm::omp::Directive::OMPD_declare_target);
+}
+
+void OmpStructureChecker::Leave(const parser::OmpDeclareTargetWithClause &x) {
+  if (x.v.v.size() > 0) {
+    const parser::OmpClause *enterClause =
+        FindClause(llvm::omp::Clause::OMPC_enter);
+    const parser::OmpClause *toClause = FindClause(llvm::omp::Clause::OMPC_to);
+    const parser::OmpClause *linkClause =
+        FindClause(llvm::omp::Clause::OMPC_link);
+    if (!enterClause && !toClause && !linkClause) {
+      context_.Say(x.source,
+          "If the DECLARE TARGET directive has a clause, it must contain at lease one ENTER clause or LINK clause"_err_en_US);
+    }
+    if (toClause) {
+      context_.Say(toClause->source,
+          "The usage of TO clause on DECLARE TARGET directive has been deprecated. Use ENTER clause instead."_warn_en_US);
+    }
+  }
+}
+
 void OmpStructureChecker::Enter(const parser::OpenMPDeclareTargetConstruct &x) {
   const auto &dir{std::get<parser::Verbatim>(x.t)};
   PushContext(dir.source, llvm::omp::Directive::OMPD_declare_target);
-  const auto &spec{std::get<parser::OmpDeclareTargetSpecifier>(x.t)};
-  if (std::holds_alternative<parser::OmpDeclareTargetWithClause>(spec.u)) {
-    SetClauseSets(llvm::omp::Directive::OMPD_declare_target);
-  }
 }
 
 void OmpStructureChecker::Enter(const parser::OmpDeclareTargetWithList &x) {
@@ -1245,7 +1263,8 @@ void OmpStructureChecker::Leave(const parser::OpenMPDeclareTargetConstruct &x) {
     CheckThreadprivateOrDeclareTargetVar(*objectList);
   } else if (const auto *clauseList{
                  parser::Unwrap<parser::OmpClauseList>(spec.u)}) {
-    bool toClauseFound{false}, deviceTypeClauseFound{false};
+    bool toClauseFound{false}, deviceTypeClauseFound{false},
+        enterClauseFound{false};
     for (const auto &clause : clauseList->v) {
       common::visit(
           common::visitors{
@@ -1260,6 +1279,12 @@ void OmpStructureChecker::Leave(const parser::OpenMPDeclareTargetConstruct &x) {
                 CheckIsVarPartOfAnotherVar(dir.source, linkClause.v);
                 CheckThreadprivateOrDeclareTargetVar(linkClause.v);
               },
+              [&](const parser::OmpClause::Enter &enterClause) {
+                enterClauseFound = true;
+                CheckSymbolNames(dir.source, enterClause.v);
+                CheckIsVarPartOfAnotherVar(dir.source, enterClause.v);
+                CheckThreadprivateOrDeclareTargetVar(enterClause.v);
+              },
               [&](const parser::OmpClause::DeviceType &deviceTypeClause) {
                 deviceTypeClauseFound = true;
                 if (deviceTypeClause.v.v !=
@@ -1273,7 +1298,7 @@ void OmpStructureChecker::Leave(const parser::OpenMPDeclareTargetConstruct &x) {
           },
           clause.u);
 
-      if (toClauseFound && !deviceTypeClauseFound) {
+      if ((toClauseFound || enterClauseFound) && !deviceTypeClauseFound) {
         deviceConstructFound_ = true;
       }
     }
@@ -2228,6 +2253,7 @@ CHECK_SIMPLE_CLAUSE(CancellationConstructType, OMPC_cancellation_construct_type)
 CHECK_SIMPLE_CLAUSE(Doacross, OMPC_doacross)
 CHECK_SIMPLE_CLAUSE(OmpxAttribute, OMPC_ompx_attribute)
 CHECK_SIMPLE_CLAUSE(OmpxBare, OMPC_ompx_bare)
+CHECK_SIMPLE_CLAUSE(Enter, OMPC_enter)
 
 CHECK_REQ_SCALAR_INT_CLAUSE(Grainsize, OMPC_grainsize)
 CHECK_REQ_SCALAR_INT_CLAUSE(NumTasks, OMPC_num_tasks)
@@ -3229,12 +3255,13 @@ const parser::OmpObjectList *OmpStructureChecker::GetOmpObjectList(
     const parser::OmpClause &clause) {
 
   // Clauses with OmpObjectList as its data member
-  using MemberObjectListClauses = std::tuple<parser::OmpClause::Copyprivate,
-      parser::OmpClause::Copyin, parser::OmpClause::Firstprivate,
-      parser::OmpClause::From, parser::OmpClause::Lastprivate,
-      parser::OmpClause::Link, parser::OmpClause::Private,
-      parser::OmpClause::Shared, parser::OmpClause::To,
-      parser::OmpClause::UseDevicePtr, parser::OmpClause::UseDeviceAddr>;
+  using MemberObjectListClauses =
+      std::tuple<parser::OmpClause::Copyprivate, parser::OmpClause::Copyin,
+          parser::OmpClause::Firstprivate, parser::OmpClause::From,
+          parser::OmpClause::Lastprivate, parser::OmpClause::Link,
+          parser::OmpClause::Private, parser::OmpClause::Shared,
+          parser::OmpClause::To, parser::OmpClause::Enter,
+          parser::OmpClause::UseDevicePtr, parser::OmpClause::UseDeviceAddr>;
 
   // Clauses with OmpObjectList in the tuple
   using TupleObjectListClauses =
diff --git a/flang/lib/Semantics/check-omp-structure.h b/flang/lib/Semantics/check-omp-structure.h
index d35602cca75d549..90e5c9f19127750 100644
--- a/flang/lib/Semantics/check-omp-structure.h
+++ b/flang/lib/Semantics/check-omp-structure.h
@@ -80,6 +80,8 @@ class OmpStructureChecker
   void Enter(const parser::OpenMPDeclareTargetConstruct &);
   void Leave(const parser::OpenMPDeclareTargetConstruct &);
   void Enter(const parser::OmpDeclareTargetWithList &);
+  void Enter(const parser::OmpDeclareTargetWithClause &);
+  void Leave(const parser::OmpDeclareTargetWithClause &);
   void Enter(const parser::OpenMPExecutableAllocate &);
   void Leave(const parser::OpenMPExecutableAllocate &);
   void Enter(const parser::OpenMPAllocatorsConstruct &);
diff --git a/flang/lib/Semantics/resolve-directives.cpp b/flang/lib/Semantics/resolve-directives.cpp
index bbb105e3516da18..882ecad1f7ed023 100644
--- a/flang/lib/Semantics/resolve-directives.cpp
+++ b/flang/lib/Semantics/resolve-directives.cpp
@@ -1744,6 +1744,9 @@ bool OmpAttributeVisitor::Pre(const parser::OpenMPDeclareTargetConstruct &x) {
       } else if (const auto *linkClause{
                      std::get_if<parser::OmpClause::Link>(&clause.u)}) {
         ResolveOmpObjectList(linkClause->v, Symbol::Flag::OmpDeclareTarget);
+      } else if (const auto *enterClause{
+                     std::get_if<parser::OmpClause::Enter>(&clause.u)}) {
+        ResolveOmpObjectList(enterClause->v, Symbol::Flag::OmpDeclareTarget);
       }
     }
   }
diff --git a/flang/test/Lower/OpenMP/FIR/declare-target-data.f90 b/flang/test/Lower/OpenMP/FIR/declare-target-data.f90
index e57d928f5497a17..bb3bbc8dfa83769 100644
--- a/flang/test/Lower/OpenMP/FIR/declare-target-data.f90
+++ b/flang/test/Lower/OpenMP/FIR/declare-target-data.f90
@@ -32,6 +32,10 @@ module test_0
 INTEGER :: data_int_to = 5
 !$omp declare target to(data_int_to)
 
+!CHECK-DAG: fir.global @_QMtest_0Edata_int_enter {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>} : i32
+INTEGER :: data_int_enter = 5
+!$omp declare target enter(data_int_enter)
+
 !CHECK-DAG: fir.global @_QMtest_0Edata_int_clauseless {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (to)>} : i32
 INTEGER :: data_int_clauseless = 1
 !$omp declare target(data_int_clauseless)
@@ -42,6 +46,12 @@ module test_0
 REAL :: data_extended_to_2 = 3
 !$omp declare target to(data_extended_to_1, data_extended_to_2)
 
+!CHECK-DAG: fir.global @_QMtest_0Edata_extended_enter_1 {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>} : f32
+!CHECK-DAG: fir.global @_QMtest_0Edata_extended_enter_2 {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>} : f32
+REAL :: data_extended_enter_1 = 2
+REAL :: data_extended_enter_2 = 3
+!$omp declare target enter(data_extended_enter_1, data_extended_enter_2)
+
 !CHECK-DAG: fir.global @_QMtest_0Edata_extended_link_1 {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (link)>} : f32
 !CHECK-DAG: fir.global @_QMtest_0Edata_extended_link_2 {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (link)>} : f32
 REAL :: data_extended_link_1 = 2
@@ -69,4 +79,10 @@ PROGRAM commons
     REAL :: two_to = 2
     COMMON /numbers_to/ one_to, two_to
     !$omp declare target to(/numbers_to/)
+
+    !CHECK-DAG: fir.global @numbers_enter_ {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>} : tuple<f32, f32> {
+    REAL :: one_enter = 1
+    REAL :: two_enter = 2
+    COMMON /numbers_enter/ one_enter, two_enter
+    !$omp declare target enter(/numbers_enter/)
 END
diff --git a/flang/test/Lower/OpenMP/FIR/declare-target-func-and-subr.f90 b/flang/test/Lower/OpenMP/FIR/declare-target-func-and-subr.f90
index 26741c6795a6ce5..36d4d7db64e5f25 100644
--- a/flang/test/Lower/OpenMP/FIR/declare-target-func-and-subr.f90
+++ b/flang/test/Lower/OpenMP/FIR/declare-target-func-and-subr.f90
@@ -13,6 +13,14 @@ FUNCTION FUNC_T_DEVICE() RESULT(I)
     I = 1
 END FUNCTION FUNC_T_DEVICE
 
+! DEVICE-LABEL: func.func @_QPfunc_enter_device()
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}
+FUNCTION FUNC_ENTER_DEVICE() RESULT(I)
+!$omp declare target enter(FUNC_ENTER_DEVICE) device_type(nohost)
+    INTEGER :: I
+    I = 1
+END FUNCTION FUNC_ENTER_DEVICE
+
 ! HOST-LABEL: func.func @_QPfunc_t_host()
 ! HOST-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (host), capture_clause = (to)>{{.*}}
 FUNCTION FUNC_T_HOST() RESULT(I)
@@ -21,6 +29,14 @@ FUNCTION FUNC_T_HOST() RESULT(I)
     I = 1
 END FUNCTION FUNC_T_HOST
 
+! HOST-LABEL: func.func @_QPfunc_enter_host()
+! HOST-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (host), capture_clause = (enter)>{{.*}}
+FUNCTION FUNC_ENTER_HOST() RESULT(I)
+!$omp declare target enter(FUNC_ENTER_HOST) device_type(host)
+    INTEGER :: I
+    I = 1
+END FUNCTION FUNC_ENTER_HOST
+
 ! ALL-LABEL: func.func @_QPfunc_t_any()
 ! ALL-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (to)>{{.*}}
 FUNCTION FUNC_T_ANY() RESULT(I)
@@ -29,6 +45,14 @@ FUNCTION FUNC_T_ANY() RESULT(I)
     I = 1
 END FUNCTION FUNC_T_ANY
 
+! ALL-LABEL: func.func @_QPfunc_enter_any()
+! ALL-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>{{.*}}
+FUNCTION FUNC_ENTER_ANY() RESULT(I)
+!$omp declare target enter(FUNC_ENTER_ANY) device_type(any)
+    INTEGER :: I
+    I = 1
+END FUNCTION FUNC_ENTER_ANY
+
 ! ALL-LABEL: func.func @_QPfunc_default_t_any()
 ! ALL-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (to)>{{.*}}
 FUNCTION FUNC_DEFAULT_T_ANY() RESULT(I)
@@ -37,6 +61,14 @@ FUNCTION FUNC_DEFAULT_T_ANY() RESULT(I)
     I = 1
 END FUNCTION FUNC_DEFAULT_T_ANY
 
+! ALL-LABEL: func.func @_QPfunc_default_enter_any()
+! ALL-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>{{.*}}
+FUNCTION FUNC_DEFAULT_ENTER_ANY() RESULT(I)
+!$omp declare target enter(FUNC_DEFAULT_ENTER_ANY)
+    INTEGER :: I
+    I = 1
+END FUNCTION FUNC_DEFAULT_ENTER_ANY
+
 ! ALL-LABEL: func.func @_QPfunc_default_any()
 ! ALL-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (to)>{{.*}}
 FUNCTION FUNC_DEFAULT_ANY() RESULT(I)
@@ -65,24 +97,48 @@ SUBROUTINE SUBR_T_DEVICE()
 !$omp declare target to(SUBR_T_DEVICE) device_type(nohost)
 END
 
+! DEVICE-LABEL: func.func @_QPsubr_enter_device()
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}
+SUBROUTINE SUBR_ENTER_DEVICE()
+!$omp declare target enter(SUBR_ENTER_DEVICE) device_type(nohost)
+END
+
 ! HOST-LABEL: func.func @_QPsubr_t_host()
 ! HOST-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (host), capture_clause = (to)>{{.*}}
 SUBROUTINE SUBR_T_HOST()
 !$omp declare target to(SUBR_T_HOST) device_type(host)
 END
 
+! HOST-LABEL: func.func @_QPsubr_enter_host()
+! HOST-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (host), capture_clause = (enter)>{{.*}}
+SUBROUTINE SUBR_ENTER_HOST()
+!$omp declare target enter(SUBR_ENTER_HOST) device_type(host)
+END
+
 ! ALL-LABEL: func.func @_QPsubr_t_any()
 ! ALL-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (to)>{{.*}}
 SUBROUTINE SUBR_T_ANY()
 !$omp declare target to(SUBR_T_ANY) device_type(any)
 END
 
+! ALL-LABEL: func.func @_QPsubr_enter_any()
+! ALL-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>{{.*}}
+SUBROUTINE SUBR_ENTER_ANY()
+!$omp declare target enter(SUBR_ENTER_ANY) device_type(any)
+END
+
 ! ALL-LABEL: func.func @_QPsubr_default_t_any()
 ! ALL-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (to)>{{.*}}
 SUBROUTINE SUBR_DEFAULT_T_ANY()
 !$omp declare target to(SUBR_DEFAULT_T_ANY)
 END
 
+! ALL-LABEL: func.func @_QPsubr_default_enter_any()
+! ALL-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>{{.*}}
+SUBROUTINE SUBR_DEFAULT_ENTER_ANY()
+!$omp declare target enter(SUBR_DEFAULT_ENTER_ANY)
+END
+
 ! ALL-LABEL: func.func @_QPsubr_default_any()
 ! ALL-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (to)>{{.*}}
 SUBROUTINE SUBR_DEFAULT_ANY()
@@ -108,3 +164,15 @@ RECURSIVE FUNCTION RECURSIVE_DECLARE_TARGET(INCREMENT) RESULT(K)
         K = RECURSIVE_DECLARE_TARGET(INCREMENT + 1)
     END IF
 END FUNCTION RECURSIVE_DECLARE_TARGET
+
+! DEVICE-LABEL: func.func @_QPrecursive_declare_target_enter
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}
+RECURSIVE FUNCTION RECURSIVE_DECLARE_TARGET_ENTER(INCREMENT) RESULT(K)
+!$omp declare target enter(RECURSIVE_DECLARE_TARGET_ENTER) device_type(nohost)
+    INTEGER :: INCREMENT, K
+    IF (INCREMENT == 10) THEN
+        K = INCREMENT
+    ELSE
+        K = RECURSIVE_DECLARE_TARGET_ENTER(INCREMENT + 1)
+    END IF
+END FUNCTION RECURSIVE_DECLARE_TARGET_ENTER
diff --git a/flang/test/Lower/OpenMP/FIR/declare-target-implicit-func-and-subr-cap-enter.f90 b/flang/test/Lower/OpenMP/FIR/declare-target-implicit-func-and-subr-cap-enter.f90
new file mode 100644
index 000000000000000..8e88d1b0f52a95f
--- /dev/null
+++ b/flang/test/Lower/OpenMP/FIR/declare-target-implicit-func-and-subr-cap-enter.f90
@@ -0,0 +1,192 @@
+!RUN: %flang_fc1 -emit-fir -fopenmp %s -o - | FileCheck %s
+!RUN: %flang_fc1 -emit-fir -fopenmp -fopenmp-is-target-device %s -o - | FileCheck %s  --check-prefix=DEVICE
+!RUN: bbc -emit-fir -fopenmp %s -o - | FileCheck %s
+!RUN: bbc -emit-fir -fopenmp -fopenmp-is-target-device %s -o - | FileCheck %s --check-prefix=DEVICE
+
+! CHECK-LABEL: func.func @_QPimplicitly_c...
[truncated]

@shraiysh
Copy link
Member Author

shraiysh commented Nov 19, 2023

One thing I wanted to point out was that with this patch clang will accept #pragma omp declare target enter(x) when compiled with version 5.2. I don't think clang is handling it accurately yet though. Is there a way to report "unimplemented" errors in clang? Or is it okay this way?

Edit: enter clause is handled in declare target in OpenMP 5.2. Just checked the generated LLVM IR and there is handling for it in ParseOpenMPDeclareTargetClauses too.

Copy link
Contributor

@skatrak skatrak left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thank you Shraiysh, I just have a couple of nits.

return findRepeatableClause<ClauseTy::Enter>(
[&](const ClauseTy::Enter *enterClause,
const Fortran::parser::CharBlock &) {
// Case: declare target to(func, var1, var2)...
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Nit: Update comment

@@ -590,6 +590,8 @@ class ClauseProcessor {
bool processSectionsReduction(mlir::Location currentLocation) const;
bool processTo(llvm::SmallVectorImpl<DeclareTargetCapturePair> &result) const;
bool
processEnter(llvm::SmallVectorImpl<DeclareTargetCapturePair> &result) const;
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Nit: It would be good to keep the alphabetic order of declarations here and definitions below, so I'd add this between processDeviceType and processFinal.

Copy link
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks for the review @skatrak. I moved it between processDepend and processIf because enter is a repeatable clause.

Copy link
Contributor

@agozillon agozillon left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM, aside from Sergio's Nits!

Might be worth asking someone from the Clang team to have a look at the Clang changes if you're a little unsure on them, but if you're happy with them can disregard doing that!

@shraiysh
Copy link
Member Author

Might be worth asking someone from the Clang team to have a look at the Clang changes if you're a little unsure on them, but if you're happy with them can disregard doing that!

Thank you for the review @agozillon @skatrak. I am not sure who I should request the review from, but I have requested a review from the codeowner of OpenMP conformance. clang:openmp is also tagged as a topic in this patch, so I will wait for 1-2 days before merging this if someone from clang/openmp team has any concerns with this.

This patch adds the following check from OpenMP 5.2.

```
If the directive has a clause, it must contain at least one enter clause
or at least one link clause.
```

Also added a warning for the deprication of `TO` clause on `DECLARE
TARGET` construct.

```
The clause-name to may be used as a synonym for the clause-name enter.
This use has been deprecated.
```
@shraiysh shraiysh merged commit 8840eb3 into llvm:main Nov 22, 2023
2 of 3 checks passed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
clang:frontend Language frontend issues, e.g. anything involving "Sema" clang:openmp OpenMP related changes to Clang clang Clang issues not falling into any other category flang:fir-hlfir flang:openmp flang:parser flang:semantics flang Flang issues not falling into any other category
Projects
None yet
Development

Successfully merging this pull request may close these issues.

4 participants