[OpenMP][Clang] Parsing/Sema support for use_device_ptr(fb_preserve/fb_nullify).

clang/include/clang/AST/OpenMPClause.h

@@ -7989,6 +7989,13 @@ class OMPUseDevicePtrClause final
   friend OMPVarListClause;
   friend TrailingObjects;
 
+  /// Fallback modifier for the clause.
+  OpenMPUseDevicePtrFallbackModifier FallbackModifier =
+      OMPC_USE_DEVICE_PTR_FALLBACK_unknown;
+
+  /// Location of the fallback modifier.
+  SourceLocation FallbackModifierLoc;
+
   /// Build clause with number of variables \a NumVars.
   ///
   /// \param Locs Locations needed to build a mappable clause. It includes 1)
@@ -7999,10 +8006,15 @@ class OMPUseDevicePtrClause final
   /// NumUniqueDeclarations: number of unique base declarations in this clause;
   /// 3) NumComponentLists: number of component lists in this clause; and 4)
   /// NumComponents: total number of expression components in the clause.
-  explicit OMPUseDevicePtrClause(const OMPVarListLocTy &Locs,
-                                 const OMPMappableExprListSizeTy &Sizes)
-      : OMPMappableExprListClause(llvm::omp::OMPC_use_device_ptr, Locs, Sizes) {
-  }
+  /// \param FallbackModifier The fallback modifier for the clause.
+  /// \param FallbackModifierLoc Location of the fallback modifier.
+  explicit OMPUseDevicePtrClause(
+      const OMPVarListLocTy &Locs, const OMPMappableExprListSizeTy &Sizes,
+      OpenMPUseDevicePtrFallbackModifier FallbackModifier,
+      SourceLocation FallbackModifierLoc)
+      : OMPMappableExprListClause(llvm::omp::OMPC_use_device_ptr, Locs, Sizes),
+        FallbackModifier(FallbackModifier),
+        FallbackModifierLoc(FallbackModifierLoc) {}
 
   /// Build an empty clause.
   ///
@@ -8055,6 +8067,14 @@ class OMPUseDevicePtrClause final
     return {getPrivateCopies().end(), varlist_size()};
   }
 
+  /// Set the fallback modifier for the clause.
+  void setFallbackModifier(OpenMPUseDevicePtrFallbackModifier M) {
+    FallbackModifier = M;
+  }
+
+  /// Set the location of the fallback modifier.
+  void setFallbackModifierLoc(SourceLocation Loc) { FallbackModifierLoc = Loc; }
+
 public:
   /// Creates clause with a list of variables \a Vars.
   ///
@@ -8067,11 +8087,15 @@ class OMPUseDevicePtrClause final
   /// \param Inits Expressions referring to private copy initializers.
   /// \param Declarations Declarations used in the clause.
   /// \param ComponentLists Component lists used in the clause.
+  /// \param FallbackModifier The fallback modifier for the clause.
+  /// \param FallbackModifierLoc Location of the fallback modifier.
   static OMPUseDevicePtrClause *
   Create(const ASTContext &C, const OMPVarListLocTy &Locs,
          ArrayRef<Expr *> Vars, ArrayRef<Expr *> PrivateVars,
          ArrayRef<Expr *> Inits, ArrayRef<ValueDecl *> Declarations,
-         MappableExprComponentListsRef ComponentLists);
+         MappableExprComponentListsRef ComponentLists,
+         OpenMPUseDevicePtrFallbackModifier FallbackModifier,
+         SourceLocation FallbackModifierLoc);
 
   /// Creates an empty clause with the place for \a NumVars variables.
   ///
@@ -8084,6 +8108,14 @@ class OMPUseDevicePtrClause final
   static OMPUseDevicePtrClause *
   CreateEmpty(const ASTContext &C, const OMPMappableExprListSizeTy &Sizes);
 
+  /// Get the fallback modifier for the clause.
+  OpenMPUseDevicePtrFallbackModifier getFallbackModifier() const {
+    return FallbackModifier;
+  }
+
+  /// Get the location of the fallback modifier.
+  SourceLocation getFallbackModifierLoc() const { return FallbackModifierLoc; }
+
   using private_copies_iterator = MutableArrayRef<Expr *>::iterator;
   using private_copies_const_iterator = ArrayRef<const Expr *>::iterator;
   using private_copies_range = llvm::iterator_range<private_copies_iterator>;

clang/include/clang/Basic/OpenMPKinds.def

@@ -110,6 +110,9 @@
 #ifndef OPENMP_NEED_DEVICE_PTR_KIND
 #define OPENMP_NEED_DEVICE_PTR_KIND(Name)
 #endif
+#ifndef OPENMP_USE_DEVICE_PTR_FALLBACK_MODIFIER
+#define OPENMP_USE_DEVICE_PTR_FALLBACK_MODIFIER(Name)
+#endif
 
 // Static attributes for 'schedule' clause.
 OPENMP_SCHEDULE_KIND(static)
@@ -282,6 +285,10 @@ OPENMP_THREADSET_KIND(omp_team)
 OPENMP_NEED_DEVICE_PTR_KIND(fb_nullify)
 OPENMP_NEED_DEVICE_PTR_KIND(fb_preserve)
 
+// OpenMP 6.1 modifiers for 'use_device_ptr' clause.
+OPENMP_USE_DEVICE_PTR_FALLBACK_MODIFIER(fb_nullify)
+OPENMP_USE_DEVICE_PTR_FALLBACK_MODIFIER(fb_preserve)
+
 #undef OPENMP_NUMTASKS_MODIFIER
 #undef OPENMP_NUMTHREADS_MODIFIER
 #undef OPENMP_DYN_GROUPPRIVATE_MODIFIER
@@ -315,3 +322,4 @@ OPENMP_NEED_DEVICE_PTR_KIND(fb_preserve)
 #undef OPENMP_ALLOCATE_MODIFIER
 #undef OPENMP_THREADSET_KIND
 #undef OPENMP_NEED_DEVICE_PTR_KIND
+#undef OPENMP_USE_DEVICE_PTR_FALLBACK_MODIFIER

clang/include/clang/Basic/OpenMPKinds.h

@@ -218,6 +218,14 @@ enum OpenMPNeedDevicePtrModifier {
   OMPC_NEED_DEVICE_PTR_unknown,
 };
 
+/// OpenMP 6.1 use_device_ptr fallback modifier
+enum OpenMPUseDevicePtrFallbackModifier {
+#define OPENMP_USE_DEVICE_PTR_FALLBACK_MODIFIER(Name)                          \
+  OMPC_USE_DEVICE_PTR_FALLBACK_##Name,
+#include "clang/Basic/OpenMPKinds.def"
+  OMPC_USE_DEVICE_PTR_FALLBACK_unknown,
+};
+
 /// OpenMP bindings for the 'bind' clause.
 enum OpenMPBindClauseKind {
 #define OPENMP_BIND_KIND(Name) OMPC_BIND_##Name,

clang/include/clang/Sema/SemaOpenMP.h

@@ -1176,6 +1176,10 @@ class SemaOpenMP : public SemaBase {
     int OriginalSharingModifier = 0; // Default is shared
     int NeedDevicePtrModifier = 0;
     SourceLocation NeedDevicePtrModifierLoc;
+    int UseDevicePtrFallbackModifier =
+        OMPC_USE_DEVICE_PTR_FALLBACK_unknown; ///< Fallback modifier for
+                                              ///< use_device_ptr clause.
+    SourceLocation UseDevicePtrFallbackModifierLoc;
     SmallVector<OpenMPMapModifierKind, NumberOfOMPMapClauseModifiers>
         MapTypeModifiers;
     SmallVector<SourceLocation, NumberOfOMPMapClauseModifiers>
@@ -1364,8 +1368,10 @@ class SemaOpenMP : public SemaBase {
                         ArrayRef<Expr *> VarList, const OMPVarListLocTy &Locs,
                         ArrayRef<Expr *> UnresolvedMappers = {});
   /// Called on well-formed 'use_device_ptr' clause.
-  OMPClause *ActOnOpenMPUseDevicePtrClause(ArrayRef<Expr *> VarList,
-                                           const OMPVarListLocTy &Locs);
+  OMPClause *ActOnOpenMPUseDevicePtrClause(
+      ArrayRef<Expr *> VarList, const OMPVarListLocTy &Locs,
+      OpenMPUseDevicePtrFallbackModifier FallbackModifier,
+      SourceLocation FallbackModifierLoc);
   /// Called on well-formed 'use_device_addr' clause.
   OMPClause *ActOnOpenMPUseDeviceAddrClause(ArrayRef<Expr *> VarList,
                                             const OMPVarListLocTy &Locs);

clang/lib/AST/OpenMPClause.cpp

@@ -1441,7 +1441,9 @@ OMPUseDevicePtrClause *OMPUseDevicePtrClause::Create(
     const ASTContext &C, const OMPVarListLocTy &Locs, ArrayRef<Expr *> Vars,
     ArrayRef<Expr *> PrivateVars, ArrayRef<Expr *> Inits,
     ArrayRef<ValueDecl *> Declarations,
-    MappableExprComponentListsRef ComponentLists) {
+    MappableExprComponentListsRef ComponentLists,
+    OpenMPUseDevicePtrFallbackModifier FallbackModifier,
+    SourceLocation FallbackModifierLoc) {
   OMPMappableExprListSizeTy Sizes;
   Sizes.NumVars = Vars.size();
   Sizes.NumUniqueDeclarations = getUniqueDeclarationsTotalNumber(Declarations);
@@ -1465,7 +1467,8 @@ OMPUseDevicePtrClause *OMPUseDevicePtrClause::Create(
           Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
           Sizes.NumComponents));
 
-  OMPUseDevicePtrClause *Clause = new (Mem) OMPUseDevicePtrClause(Locs, Sizes);
+  OMPUseDevicePtrClause *Clause = new (Mem)
+      OMPUseDevicePtrClause(Locs, Sizes, FallbackModifier, FallbackModifierLoc);
 
   Clause->setVarRefs(Vars);
   Clause->setPrivateCopies(PrivateVars);
@@ -2753,7 +2756,15 @@ void OMPClausePrinter::VisitOMPDefaultmapClause(OMPDefaultmapClause *Node) {
 void OMPClausePrinter::VisitOMPUseDevicePtrClause(OMPUseDevicePtrClause *Node) {
   if (!Node->varlist_empty()) {
     OS << "use_device_ptr";
-    VisitOMPClauseList(Node, '(');
+    if (Node->getFallbackModifier() != OMPC_USE_DEVICE_PTR_FALLBACK_unknown) {
+      OS << "("
+         << getOpenMPSimpleClauseTypeName(OMPC_use_device_ptr,
+                                          Node->getFallbackModifier())
+         << ":";
+      VisitOMPClauseList(Node, ' ');
+    } else {
+      VisitOMPClauseList(Node, '(');
+    }
     OS << ")";
   }
 }

clang/lib/Basic/OpenMPKinds.cpp

@@ -238,6 +238,16 @@ unsigned clang::getOpenMPSimpleClauseType(OpenMPClauseKind Kind, StringRef Str,
       return OMPC_NUMTHREADS_unknown;
     return Type;
   }
+  case OMPC_use_device_ptr: {
+    unsigned Type = llvm::StringSwitch<unsigned>(Str)
+#define OPENMP_USE_DEVICE_PTR_FALLBACK_MODIFIER(Name)                          \
+  .Case(#Name, OMPC_USE_DEVICE_PTR_FALLBACK_##Name)
+#include "clang/Basic/OpenMPKinds.def"
+                        .Default(OMPC_USE_DEVICE_PTR_FALLBACK_unknown);
+    if (LangOpts.OpenMP < 61)
+      return OMPC_USE_DEVICE_PTR_FALLBACK_unknown;
+    return Type;
+  }
   case OMPC_unknown:
   case OMPC_threadprivate:
   case OMPC_groupprivate:
@@ -280,7 +290,6 @@ unsigned clang::getOpenMPSimpleClauseType(OpenMPClauseKind Kind, StringRef Str,
   case OMPC_nogroup:
   case OMPC_hint:
   case OMPC_uniform:
-  case OMPC_use_device_ptr:
   case OMPC_use_device_addr:
   case OMPC_is_device_ptr:
   case OMPC_has_device_addr:
@@ -608,6 +617,16 @@ const char *clang::getOpenMPSimpleClauseTypeName(OpenMPClauseKind Kind,
 #include "clang/Basic/OpenMPKinds.def"
     }
     llvm_unreachable("Invalid OpenMP 'threadset' clause modifier");
+  case OMPC_use_device_ptr:
+    switch (Type) {
+    case OMPC_USE_DEVICE_PTR_FALLBACK_unknown:
+      return "unknown";
+#define OPENMP_USE_DEVICE_PTR_FALLBACK_MODIFIER(Name)                          \
+  case OMPC_USE_DEVICE_PTR_FALLBACK_##Name:                                    \
+    return #Name;
+#include "clang/Basic/OpenMPKinds.def"
+    }
+    llvm_unreachable("Invalid OpenMP 'use_device_ptr' clause modifier");
   case OMPC_unknown:
   case OMPC_threadprivate:
   case OMPC_groupprivate:
@@ -650,7 +669,6 @@ const char *clang::getOpenMPSimpleClauseTypeName(OpenMPClauseKind Kind,
   case OMPC_nogroup:
   case OMPC_hint:
   case OMPC_uniform:
-  case OMPC_use_device_ptr:
   case OMPC_use_device_addr:
   case OMPC_is_device_ptr:
   case OMPC_has_device_addr:

clang/lib/Parse/ParseOpenMP.cpp

@@ -5053,6 +5053,23 @@ bool Parser::ParseOpenMPVarList(OpenMPDirectiveKind DKind,
       ExpectAndConsume(tok::colon, diag::warn_pragma_expected_colon,
                        "adjust-op");
     }
+  } else if (Kind == OMPC_use_device_ptr && getLangOpts().OpenMP >= 61) {
+    // Handle optional fallback modifier for use_device_ptr clause.
+    // use_device_ptr([fb_preserve | fb_nullify :] list)
+    // Default is fb_preserve.
+    if (Tok.is(tok::identifier)) {
+      auto FallbackModifier = static_cast<OpenMPUseDevicePtrFallbackModifier>(
+          getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok), getLangOpts()));
+      if (FallbackModifier != OMPC_USE_DEVICE_PTR_FALLBACK_unknown) {
+        Data.UseDevicePtrFallbackModifier = FallbackModifier;
+        Data.UseDevicePtrFallbackModifierLoc = Tok.getLocation();
+        ConsumeToken();
+        if (Tok.is(tok::colon))
+          Data.ColonLoc = ConsumeToken();
+        else
+          Diag(Tok, diag::err_modifier_expected_colon) << "fallback";
+      }
+    }
   }
 
   bool IsComma =

clang/lib/Sema/SemaOpenMP.cpp

@@ -18724,7 +18724,11 @@ OMPClause *SemaOpenMP::ActOnOpenMPVarListClause(OpenMPClauseKind Kind,
         VarList, Locs);
     break;
   case OMPC_use_device_ptr:
-    Res = ActOnOpenMPUseDevicePtrClause(VarList, Locs);
+    Res = ActOnOpenMPUseDevicePtrClause(
+        VarList, Locs,
+        static_cast<OpenMPUseDevicePtrFallbackModifier>(
+            Data.UseDevicePtrFallbackModifier),
+        Data.UseDevicePtrFallbackModifierLoc);
     break;
   case OMPC_use_device_addr:
     Res = ActOnOpenMPUseDeviceAddrClause(VarList, Locs);
@@ -24539,9 +24543,10 @@ OMPClause *SemaOpenMP::ActOnOpenMPFromClause(
       MapperId);
 }
 
-OMPClause *
-SemaOpenMP::ActOnOpenMPUseDevicePtrClause(ArrayRef<Expr *> VarList,
-                                          const OMPVarListLocTy &Locs) {
+OMPClause *SemaOpenMP::ActOnOpenMPUseDevicePtrClause(
+    ArrayRef<Expr *> VarList, const OMPVarListLocTy &Locs,
+    OpenMPUseDevicePtrFallbackModifier FallbackModifier,
+    SourceLocation FallbackModifierLoc) {
   MappableVarListInfo MVLI(VarList);
   SmallVector<Expr *, 8> PrivateCopies;
   SmallVector<Expr *, 8> Inits;
@@ -24622,7 +24627,8 @@ SemaOpenMP::ActOnOpenMPUseDevicePtrClause(ArrayRef<Expr *> VarList,
 
   return OMPUseDevicePtrClause::Create(
       getASTContext(), Locs, MVLI.ProcessedVarList, PrivateCopies, Inits,
-      MVLI.VarBaseDeclarations, MVLI.VarComponents);
+      MVLI.VarBaseDeclarations, MVLI.VarComponents, FallbackModifier,
+      FallbackModifierLoc);
 }
 
 OMPClause *

clang/lib/Sema/TreeTransform.h

@@ -2252,9 +2252,12 @@ class TreeTransform {
   ///
   /// By default, performs semantic analysis to build the new OpenMP clause.
   /// Subclasses may override this routine to provide different behavior.
-  OMPClause *RebuildOMPUseDevicePtrClause(ArrayRef<Expr *> VarList,
-                                          const OMPVarListLocTy &Locs) {
-    return getSema().OpenMP().ActOnOpenMPUseDevicePtrClause(VarList, Locs);
+  OMPClause *RebuildOMPUseDevicePtrClause(
+      ArrayRef<Expr *> VarList, const OMPVarListLocTy &Locs,
+      OpenMPUseDevicePtrFallbackModifier FallbackModifier,
+      SourceLocation FallbackModifierLoc) {
+    return getSema().OpenMP().ActOnOpenMPUseDevicePtrClause(
+        VarList, Locs, FallbackModifier, FallbackModifierLoc);
   }
 
   /// Build a new OpenMP 'use_device_addr' clause.
@@ -11591,7 +11594,8 @@ OMPClause *TreeTransform<Derived>::TransformOMPUseDevicePtrClause(
     Vars.push_back(EVar.get());
   }
   OMPVarListLocTy Locs(C->getBeginLoc(), C->getLParenLoc(), C->getEndLoc());
-  return getDerived().RebuildOMPUseDevicePtrClause(Vars, Locs);
+  return getDerived().RebuildOMPUseDevicePtrClause(
+      Vars, Locs, C->getFallbackModifier(), C->getFallbackModifierLoc());
 }
 
 template <typename Derived>

clang/lib/Serialization/ASTReader.cpp

@@ -12503,6 +12503,8 @@ void OMPClauseReader::VisitOMPFromClause(OMPFromClause *C) {
 
 void OMPClauseReader::VisitOMPUseDevicePtrClause(OMPUseDevicePtrClause *C) {
   C->setLParenLoc(Record.readSourceLocation());
+  C->setFallbackModifier(Record.readEnum<OpenMPUseDevicePtrFallbackModifier>());
+  C->setFallbackModifierLoc(Record.readSourceLocation());
   auto NumVars = C->varlist_size();
   auto UniqueDecls = C->getUniqueDeclarationsNum();
   auto TotalLists = C->getTotalComponentListNum();

clang/lib/Serialization/ASTWriter.cpp

@@ -8487,6 +8487,8 @@ void OMPClauseWriter::VisitOMPUseDevicePtrClause(OMPUseDevicePtrClause *C) {
   Record.push_back(C->getTotalComponentListNum());
   Record.push_back(C->getTotalComponentsNum());
   Record.AddSourceLocation(C->getLParenLoc());
+  Record.writeEnum(C->getFallbackModifier());
+  Record.AddSourceLocation(C->getFallbackModifierLoc());
   for (auto *E : C->varlist())
     Record.AddStmt(E);
   for (auto *VE : C->private_copies())

clang/test/OpenMP/target_data_use_device_ptr_fallback_ast_print.cpp

@@ -0,0 +1,36 @@
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=61 -ast-print %s | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=61 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=61 -std=c++11 -include-pch %t -verify %s -ast-print | FileCheck %s
+
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+// CHECK-LABEL:void f1(int *p, int *q)
+void f1(int *p, int *q) {
+
+// CHECK: #pragma omp target data use_device_ptr(fb_preserve: p)
+#pragma omp target data use_device_ptr(fb_preserve: p)
+  {}
+
+// CHECK: #pragma omp target data use_device_ptr(fb_nullify: p)
+#pragma omp target data use_device_ptr(fb_nullify: p)
+  {}
+
+// Without any fallback modifier
+// CHECK: #pragma omp target data use_device_ptr(p)
+#pragma omp target data use_device_ptr(p)
+  {}
+
+// Multiple variables with fb_preserve
+// CHECK: #pragma omp target data use_device_ptr(fb_preserve: p,q)
+#pragma omp target data use_device_ptr(fb_preserve: p, q)
+  {}
+
+// Multiple variables with fb_nullify
+// CHECK: #pragma omp target data use_device_ptr(fb_nullify: p,q)
+#pragma omp target data use_device_ptr(fb_nullify: p, q)
+  {}
+}
+#endif