[OpenMP] Initial implementation of parse+sema for clause use_device_ptr of 'target data'

[carlo.bertolli]

Summary

This patch is similar to the implementation of 'private' clause: it adds a list of private pointers to be used within the target data region to store the device pointers returned by the runtime.
Please refer to the following document for a full description of what the runtime witll return in this case (page 10 and 11):
https://github.com/clang-omp/OffloadingDesign

I am happy to answer any question related to the runtime interface to help reviewing this patch.

[carlo.bertolli]

[OpenMP] Delete trailing whitespaces in regression test file and rebase.

[ABataev]

include/clang/Basic/DiagnosticSemaKinds.td

8318

It's better to rephrase this message: "expected a variable of pointer type" or something like this. Also, it would be good to have a corresponding note, describing the provided item.

lib/Sema/SemaOpenMP.cpp

11377

You should not skip the reference from the type, references are not allowed also

11382–11394

I don't think that this patch requires some private copies right now. They will be required only during codegen. I think the private copies must be excluded from parsing/sema patch

[sfantao]

lib/Sema/SemaOpenMP.cpp

11377

Hi Alexey, I think the item can be a reference to a pointer. I couldn't find that restriction in the spec, can you point me to what you are referring to? We should probably refer to that in the preceding comment.

Thanks!

[ABataev]

lib/Sema/SemaOpenMP.cpp

11377

Here is the restriction:
"2.10.1 target data Construct, restrictions. References in the construct to a list item that appears in a use_device_ptr clause must be to the address of the list item."
Does it mean that references are allowed?

[sfantao]

lib/Sema/SemaOpenMP.cpp

11377

Ok, thanks. My interpretation is that that restriction refers to references used in clauses (in the same construct) other than use_device_ptr to declarations used in use_device_ptr. I think the motivation is to restrict the use of the value of the pointer by other clauses given that value is going to change for that data environment because of the behavior implemented by use_device_ptr. Carlo, what is your interpretation?

[kkwli0]

lib/Sema/SemaOpenMP.cpp

11377

I interpret the rule to the list item being used inside the construct instead of being used in the other clauses on the same construct.

[carlo.bertolli]

I think that Kelvin is right. This is what that sentence refers to:
int * bla = ..;
#pragma omp target data use_device_ptr(bla)
{

.. bla .. // <-- this is the 'reference' that sentence is about

}

What puzzles me about that sentence is that it seems it is asking us to use bla as following:
int addr_bla = &bla;
// now can use addr_bla

I am specifically referring to this part of the sentence: "must be to the address of the list item"

Anyway, as Kelvin says, this has nothing to do with what we expect sema for use_device_ptr to accept. Samuel point, but I may be mistaken, is that a reference to a pointer should be considered as a pointer itself.
Alexey: if you do not think this is right, I can add a test that excludes references.

Finally, thanks for the very quick review and for all your comments

• Carlo

[ABataev]

In D21904#472074, @carlo.bertolli wrote:

I think that Kelvin is right. This is what that sentence refers to:
int * bla = ..;
#pragma omp target data use_device_ptr(bla)
{

.. bla .. // <-- this is the 'reference' that sentence is about

}

What puzzles me about that sentence is that it seems it is asking us to use bla as following:
int addr_bla = &bla;
// now can use addr_bla

I am specifically referring to this part of the sentence: "must be to the address of the list item"

Anyway, as Kelvin says, this has nothing to do with what we expect sema for use_device_ptr to accept. Samuel point, but I may be mistaken, is that a reference to a pointer should be considered as a pointer itself.
Alexey: if you do not think this is right, I can add a test that excludes references.

Finally, thanks for the very quick review and for all your comments

• Carlo

Carlo, I'm not sure about this sentence at all. If you think that this construct may accept references to pointers, then go ahead and ignore my previous comment.

[carlo.bertolli]

[OpenMP] Remove private variable creation from this patch as requested by comments and rebase.

[ABataev]

include/clang/AST/OpenMPClause.h

4221

No \brief's

include/clang/Basic/DiagnosticSemaKinds.td

8333

I think it 's better to use 'expected ...' form as in other errors/warnings

include/clang/Sema/Sema.h

8466

No \brief

lib/Sema/TreeTransform.h

1784

Remove \brief tag

[carlo.bertolli]

[OpenMP] Remove 'brief' comment tags and change error report to standard 'expected' form.

[ABataev]

No positive tests for the construct

include/clang/Basic/DiagnosticSemaKinds.td

8333

Shall we accept refs to pointers and arrays?

lib/Sema/SemaOpenMP.cpp

11644–11646

It is better to add this code in codegen part

[carlo.bertolli]

No positive tests for the construct

Do you mean in the regression test? I thought I covered all negative and positive cases, but please let me know if I missed anything.

include/clang/Basic/DiagnosticSemaKinds.td

8333

I think that people of the OpenMP committee agreed that we can accept references to pointers. There is no indication in the specifications about arrays, so I will leave them out for the time being.
I will add a mention of references here.

[carlo.bertolli]

[OpenMP] Apply comments: remove creation of private variable and update regression tests to include references to pointers.

[ABataev]

I don't see successful test for the clause, which prints AST and checks serialization/deserialization. Also it should include a template test to check that dependent objects are handled correctly

include/clang/AST/OpenMPClause.h

4236

\brief

4258

\brief

4269

\brief

lib/Sema/SemaOpenMP.cpp

11648

What if Type is dependent?

[carlo.bertolli]

lib/Sema/SemaOpenMP.cpp

11648

In that case at line 8084 (function getPrivateItem) will return the pair (nullptr, true) and at line 11640 I will push nullptr for later analysis.
If this is not enough, can you please explain what needs to be done? Thanks!!

[carlo.bertolli]

[OpenMP] Add regression test to check correctness of ast building, remove 'brief's, and rebase.

[ABataev]

LG

[carlo.bertolli]

Committed revision 275271.

include/clang/AST/OpenMPClause.h

                        reinterpret_cast<Stmt **>(varlist_end()));
   }
 };
+
+/// This represents clause 'use_device_ptr' in the '#pragma omp ...'
+/// directives.
+///
+/// \code
+/// #pragma omp target data use_device_ptr(a,b)
+/// \endcode
+/// In this example directive '#pragma omp target data' has clause
+/// 'use_device_ptr' with the variables 'a' and 'b'.
+///
+class OMPUseDevicePtrClause final
+    : public OMPVarListClause<OMPUseDevicePtrClause>,
+      private llvm::TrailingObjects<OMPUseDevicePtrClause, Expr *> {
+  friend TrailingObjects;
+  friend OMPVarListClause;
+  friend class OMPClauseReader;
+  /// \brief Build clause with number of variables \a N.
+  ///
+  /// \param StartLoc Starting location of the clause.
+  /// \param LParenLoc Location of '('.
+  /// \param EndLoc Ending location of the clause.
+  /// \param N Number of the variables in the clause.
+  ///
+  OMPUseDevicePtrClause(SourceLocation StartLoc, SourceLocation LParenLoc,
+                        SourceLocation EndLoc, unsigned N)
+      : OMPVarListClause<OMPUseDevicePtrClause>(OMPC_use_device_ptr, StartLoc,
+                                                LParenLoc, EndLoc, N) {}
+
+  /// \brief Build an empty clause.
+  ///
+  /// \param N Number of variables.
+  ///
+  explicit OMPUseDevicePtrClause(unsigned N)
+      : OMPVarListClause<OMPUseDevicePtrClause>(
+            OMPC_use_device_ptr, SourceLocation(), SourceLocation(),
+            SourceLocation(), N) {}
+
+public:
+  /// \brief Creates clause with a list of variables \a VL.
+  ///
+  /// \param C AST context.
+  /// \param StartLoc Starting location of the clause.
+  /// \param LParenLoc Location of '('.
+  /// \param EndLoc Ending location of the clause.
+  /// \param VL List of references to the variables.
+  ///
+  static OMPUseDevicePtrClause *
+  Create(const ASTContext &C, SourceLocation StartLoc, SourceLocation LParenLoc,
+         SourceLocation EndLoc, ArrayRef<Expr *> VL);
+  /// \brief Creates an empty clause with the place for \a N variables.
+  ///
+  /// \param C AST context.
+  /// \param N The number of variables.
+  ///
+  static OMPUseDevicePtrClause *CreateEmpty(const ASTContext &C, unsigned N);
+
+  child_range children() {
+    return child_range(reinterpret_cast<Stmt **>(varlist_begin()),
+                       reinterpret_cast<Stmt **>(varlist_end()));
+  }
+
+  static bool classof(const OMPClause *T) {
+    return T->getClauseKind() == OMPC_use_device_ptr;
+  }
+};
 } // end namespace clang
 
 #endif // LLVM_CLANG_AST_OPENMPCLAUSE_H

include/clang/AST/RecursiveASTVisitor.h

   return true;
 }
 
+template <typename Derived>
+bool RecursiveASTVisitor<Derived>::VisitOMPUseDevicePtrClause(
+    OMPUseDevicePtrClause *C) {
+  TRY_TO(VisitOMPClauseList(C));
+  return true;
+}
+
 // FIXME: look at the following tricky-seeming exprs to see if we
 // need to recurse on anything.  These are ones that have methods
 // returning decls or qualtypes or nestednamespecifier -- though I'm

include/clang/Basic/DiagnosticSemaKinds.td

   "expected a reference to an integer-typed parameter">;
 def err_omp_at_least_one_motion_clause_required : Error<
   "expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'">;
+def  err_omp_usedeviceptr_not_a_pointer : Error<
+  "expected pointer or reference to pointer in 'use_device_ptr' clause">;
 def warn_omp_nesting_simd : Warning<
   "OpenMP only allows an ordered construct with the simd clause nested in a simd construct">,
   InGroup<SourceUsesOpenMP>;

include/clang/Basic/OpenMPKinds.def

 OPENMP_CLAUSE(defaultmap, OMPDefaultmapClause)
 OPENMP_CLAUSE(to, OMPToClause)
 OPENMP_CLAUSE(from, OMPFromClause)
+OPENMP_CLAUSE(use_device_ptr, OMPUseDevicePtrClause)
 
 // Clauses allowed for OpenMP directive 'parallel'.
 OPENMP_PARALLEL_CLAUSE(if)

 OPENMP_TARGET_DATA_CLAUSE(if)
 OPENMP_TARGET_DATA_CLAUSE(device)
 OPENMP_TARGET_DATA_CLAUSE(map)
+OPENMP_TARGET_DATA_CLAUSE(use_device_ptr)
 
 // Clauses allowed for OpenMP directive 'target enter data'.
 OPENMP_TARGET_ENTER_DATA_CLAUSE(if)

include/clang/Sema/Sema.h

                                    SourceLocation StartLoc,
                                    SourceLocation LParenLoc,
                                    SourceLocation EndLoc);
+  /// Called on well-formed 'use_device_ptr' clause.
+  OMPClause *ActOnOpenMPUseDevicePtrClause(ArrayRef<Expr *> VarList,
+                                           SourceLocation StartLoc,
+                                           SourceLocation LParenLoc,
+                                           SourceLocation EndLoc);
 
   /// \brief The kind of conversion being performed.
   enum CheckedConversionKind {

lib/AST/OpenMPClause.cpp

   case OMPC_uniform:
   case OMPC_to:
   case OMPC_from:
+  case OMPC_use_device_ptr:
     break;
   }
 

   case OMPC_uniform:
   case OMPC_to:
   case OMPC_from:
+  case OMPC_use_device_ptr:
     break;
   }
 

   return new (Mem) OMPFromClause(NumVars, NumUniqueDeclarations,
                                  NumComponentLists, NumComponents);
 }
+
+OMPUseDevicePtrClause *OMPUseDevicePtrClause::Create(const ASTContext &C,
+                                                     SourceLocation StartLoc,
+                                                     SourceLocation LParenLoc,
+                                                     SourceLocation EndLoc,
+                                                     ArrayRef<Expr *> VL) {
+  void *Mem = C.Allocate(totalSizeToAlloc<Expr *>(VL.size()));
+  OMPUseDevicePtrClause *Clause =
+      new (Mem) OMPUseDevicePtrClause(StartLoc, LParenLoc, EndLoc, VL.size());
+  Clause->setVarRefs(VL);
+  return Clause;
+}
+
+OMPUseDevicePtrClause *OMPUseDevicePtrClause::CreateEmpty(const ASTContext &C,
+                                                          unsigned N) {
+  void *Mem = C.Allocate(totalSizeToAlloc<Expr *>(N));
+  return new (Mem) OMPUseDevicePtrClause(N);
+}

lib/AST/StmtPrinter.cpp

     Node->getDefaultmapKind());
   OS << ")";
 }
+
+void OMPClausePrinter::VisitOMPUseDevicePtrClause(OMPUseDevicePtrClause *Node) {
+  if (!Node->varlist_empty()) {
+    OS << "use_device_ptr";
+    VisitOMPClauseList(Node, '(');
+    OS << ")";
+  }
+}
 }
 
 //===----------------------------------------------------------------------===//

lib/AST/StmtProfile.cpp

 void OMPClauseProfiler::VisitOMPFromClause(const OMPFromClause *C) {
   VisitOMPClauseList(C);
 }
+void OMPClauseProfiler::VisitOMPUseDevicePtrClause(
+    const OMPUseDevicePtrClause *C) {
+  VisitOMPClauseList(C);
+}
 }
 
 void

lib/Basic/OpenMPKinds.cpp

   case OMPC_uniform:
   case OMPC_to:
   case OMPC_from:
+  case OMPC_use_device_ptr:
     break;
   }
   llvm_unreachable("Invalid OpenMP simple clause kind");

   case OMPC_uniform:
   case OMPC_to:
   case OMPC_from:
+  case OMPC_use_device_ptr:
     break;
   }
   llvm_unreachable("Invalid OpenMP simple clause kind");

lib/CodeGen/CGStmtOpenMP.cpp

   case OMPC_uniform:
   case OMPC_to:
   case OMPC_from:
+  case OMPC_use_device_ptr:
     llvm_unreachable("Clause is not allowed in 'omp atomic'.");
   }
 }

lib/Parse/ParseOpenMP.cpp

   case OMPC_map:
   case OMPC_to:
   case OMPC_from:
+  case OMPC_use_device_ptr:
     Clause = ParseOpenMPVarListClause(DKind, CKind);
     break;
   case OMPC_unknown:

 ///       'to' '(' list ')'
 ///    from-clause:
 ///       'from' '(' list ')'
+///    use_device_ptr-clause:
+///       'use_device_ptr' '(' list ')'
 ///
 /// For 'linear' clause linear-list may have the following forms:
 ///  list

lib/Sema/SemaOpenMP.cpp

   case OMPC_uniform:
   case OMPC_to:
   case OMPC_from:
+  case OMPC_use_device_ptr:
     llvm_unreachable("Clause is not allowed.");
   }
   return Res;

   case OMPC_uniform:
   case OMPC_to:
   case OMPC_from:
+  case OMPC_use_device_ptr:
     llvm_unreachable("Clause is not allowed.");
   }
   return Res;

   case OMPC_uniform:
   case OMPC_to:
   case OMPC_from:
+  case OMPC_use_device_ptr:
     llvm_unreachable("Clause is not allowed.");
   }
   return Res;

   case OMPC_uniform:
   case OMPC_to:
   case OMPC_from:
+  case OMPC_use_device_ptr:
     llvm_unreachable("Clause is not allowed.");
   }
   return Res;

   case OMPC_from:
     Res = ActOnOpenMPFromClause(VarList, StartLoc, LParenLoc, EndLoc);
     break;
+  case OMPC_use_device_ptr:
+    Res = ActOnOpenMPUseDevicePtrClause(VarList, StartLoc, LParenLoc, EndLoc);
+    break;
   case OMPC_if:
   case OMPC_final:
   case OMPC_num_threads:

                                MVLI.ProcessedVarList, MVLI.VarBaseDeclarations,
                                MVLI.VarComponents);
 }
+
+OMPClause *Sema::ActOnOpenMPUseDevicePtrClause(ArrayRef<Expr *> VarList,
+                                               SourceLocation StartLoc,
+                                               SourceLocation LParenLoc,
+                                               SourceLocation EndLoc) {
+  SmallVector<Expr *, 8> Vars;
+  for (auto &RefExpr : VarList) {
+    assert(RefExpr && "NULL expr in OpenMP use_device_ptr clause.");
+    SourceLocation ELoc;
+    SourceRange ERange;
+    Expr *SimpleRefExpr = RefExpr;
+    auto Res = getPrivateItem(*this, SimpleRefExpr, ELoc, ERange);
+    if (Res.second) {
+      // It will be analyzed later.
+      Vars.push_back(RefExpr);
+    }
+    ValueDecl *D = Res.first;
+    if (!D)
+      continue;
+
+    QualType Type = D->getType();
+    // item should be a pointer or reference to pointer
+    if (!Type.getNonReferenceType()->isPointerType()) {
+      Diag(ELoc, diag::err_omp_usedeviceptr_not_a_pointer)
+          << 0 << RefExpr->getSourceRange();
+      continue;
+    }
+    Vars.push_back(RefExpr->IgnoreParens());
+  }
+
+  if (Vars.empty())
+    return nullptr;
+
+  return OMPUseDevicePtrClause::Create(Context, StartLoc, LParenLoc, EndLoc,
+                                       Vars);
+}

lib/Sema/TreeTransform.h

                                            EndLoc);
   }
 
+  /// Build a new OpenMP 'use_device_ptr' clause.
+  ///
+  /// 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,
+                                          SourceLocation StartLoc,
+                                          SourceLocation LParenLoc,
+                                          SourceLocation EndLoc) {
+    return getSema().ActOnOpenMPUseDevicePtrClause(VarList, StartLoc, LParenLoc,
+                                                   EndLoc);
+  }
+
   /// \brief Rebuild the operand to an Objective-C \@synchronized statement.
   ///
   /// By default, performs semantic analysis to build the new statement.

                                            C->getLParenLoc(), C->getLocEnd());
 }
 
+template <typename Derived>
+OMPClause *TreeTransform<Derived>::TransformOMPUseDevicePtrClause(
+    OMPUseDevicePtrClause *C) {
+  llvm::SmallVector<Expr *, 16> Vars;
+  Vars.reserve(C->varlist_size());
+  for (auto *VE : C->varlists()) {
+    ExprResult EVar = getDerived().TransformExpr(cast<Expr>(VE));
+    if (EVar.isInvalid())
+      return nullptr;
+    Vars.push_back(EVar.get());
+  }
+  return getDerived().RebuildOMPUseDevicePtrClause(
+      Vars, C->getLocStart(), C->getLParenLoc(), C->getLocEnd());
+}
+
 //===----------------------------------------------------------------------===//
 // Expression transformation
 //===----------------------------------------------------------------------===//

lib/Serialization/ASTReaderStmt.cpp

                                    NumComponents);
     break;
   }
+  case OMPC_use_device_ptr:
+    C = OMPUseDevicePtrClause::CreateEmpty(Context, Record[Idx++]);
+    break;
   }
   Visit(C);
   C->setLocStart(Reader->ReadSourceLocation(Record, Idx));

   C->setComponents(Components, ListSizes);
 }
 
+void OMPClauseReader::VisitOMPUseDevicePtrClause(OMPUseDevicePtrClause *C) {
+  C->setLParenLoc(Reader->ReadSourceLocation(Record, Idx));
+  unsigned NumVars = C->varlist_size();
+  SmallVector<Expr *, 16> Vars;
+  Vars.reserve(NumVars);
+  for (unsigned i = 0; i != NumVars; ++i)
+    Vars.push_back(Reader->Reader.ReadSubExpr());
+  C->setVarRefs(Vars);
+  Vars.clear();
+}
+
 //===----------------------------------------------------------------------===//
 // OpenMP Directives.
 //===----------------------------------------------------------------------===//

lib/Serialization/ASTWriterStmt.cpp

   }
 }
 
+void OMPClauseWriter::VisitOMPUseDevicePtrClause(OMPUseDevicePtrClause *C) {
+  Record.push_back(C->varlist_size());
+  Record.AddSourceLocation(C->getLParenLoc());
+  for (auto *VE : C->varlists()) {
+    Record.AddStmt(VE);
+  }
+}
+
 //===----------------------------------------------------------------------===//
 // OpenMP Directives.
 //===----------------------------------------------------------------------===//

test/OpenMP/target_data_use_device_ptr_messages.cpp

+// RUN: %clang_cc1 -std=c++11 -verify -fopenmp -ferror-limit 200 %s
+struct ST {
+  int *a;
+};
+struct SA {
+  const int d = 5;
+  const int da[5] = { 0 };
+  ST e;
+  ST g[10];
+  int i;
+  int &j = i;
+  int *k = &j;
+  int *&z = k;
+  int aa[10];
+  void func(int arg) {
+#pragma omp target data map(i) use_device_ptr // expected-error {{expected '(' after 'use_device_ptr'}}
+    {}
+#pragma omp target data map(i) use_device_ptr( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
+    {}
+#pragma omp target data map(i) use_device_ptr() // expected-error {{expected expression}}
+    {}
+#pragma omp target data map(i) use_device_ptr(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
+    {}
+#pragma omp target data map(i) use_device_ptr(arg // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+    {}
+#pragma omp target data map(i) use_device_ptr(i) // expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+    {}
+#pragma omp target data map(i) use_device_ptr(j) // expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+    {}
+#pragma omp target data map(i) use_device_ptr(k) // OK
+    {}
+#pragma omp target data map(i) use_device_ptr(z) // OK
+    {}
+#pragma omp target data map(i) use_device_ptr(aa) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+    {}
+#pragma omp target data map(i) use_device_ptr(e) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+    {}
+#pragma omp target data map(i) use_device_ptr(g) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+    {}
+#pragma omp target data map(i) use_device_ptr(k,i,j) // expected-error2 {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+    {}
+#pragma omp target data map(i) use_device_ptr(d) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+    {}
+#pragma omp target data map(i) use_device_ptr(da) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+    {}
+  return;
+ }
+};
+struct SB {
+  unsigned A;
+  unsigned B;
+  float Arr[100];
+  float *Ptr;
+  float *foo() {
+    return &Arr[0];
+  }
+};
+
+struct SC {
+  unsigned A : 2;
+  unsigned B : 3;
+  unsigned C;
+  unsigned D;
+  float Arr[100];
+  SB S;
+  SB ArrS[100];
+  SB *PtrS;
+  SB *&RPtrS;
+  float *Ptr;
+
+  SC(SB *&_RPtrS) : RPtrS(_RPtrS) {}
+};
+
+union SD {
+  unsigned A;
+  float B;
+};
+
+struct S1;
+extern S1 a;
+class S2 {
+  mutable int a;
+public:
+  S2():a(0) { }
+  S2(S2 &s2):a(s2.a) { }
+  static float S2s;
+  static const float S2sc;
+};
+const float S2::S2sc = 0;
+const S2 b;
+const S2 ba[5];
+class S3 {
+  int a;
+public:
+  S3():a(0) { }
+  S3(S3 &s3):a(s3.a) { }
+};
+const S3 c;
+const S3 ca[5];
+extern const int f;
+class S4 {
+  int a;
+  S4();
+  S4(const S4 &s4);
+public:
+  S4(int v):a(v) { }
+};
+class S5 {
+  int a;
+  S5():a(0) {}
+  S5(const S5 &s5):a(s5.a) { }
+public:
+  S5(int v):a(v) { }
+};
+
+S3 h;
+#pragma omp threadprivate(h)
+
+typedef int from;
+
+template <typename T, int I>
+T tmain(T argc) {
+  const T d = 5;
+  const T da[5] = { 0 };
+  S4 e(4);
+  S5 g(5);
+  T i;
+  T &j = i;
+  T *k = &j;
+  T *&z = k;
+  T aa[10];
+#pragma omp target data map(i) use_device_ptr // expected-error {{expected '(' after 'use_device_ptr'}}
+  {}
+#pragma omp target data map(i) use_device_ptr( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
+  {}
+#pragma omp target data map(i) use_device_ptr() // expected-error {{expected expression}}
+  {}
+#pragma omp target data map(i) use_device_ptr(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
+  {}
+#pragma omp target data map(i) use_device_ptr(argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_ptr(i) // expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_ptr(j) // expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_ptr(k) // OK
+  {}
+#pragma omp target data map(i) use_device_ptr(z) // OK
+  {}
+#pragma omp target data map(i) use_device_ptr(aa) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_ptr(e) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_ptr(g) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_ptr(k,i,j) // expected-error2 {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_ptr(d) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_ptr(da) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+  return 0;
+}
+
+int main(int argc, char **argv) {
+  const int d = 5;
+  const int da[5] = { 0 };
+  S4 e(4);
+  S5 g(5);
+  int i;
+  int &j = i;
+  int *k = &j;
+  int *&z = k;
+  int aa[10];
+#pragma omp target data map(i) use_device_ptr // expected-error {{expected '(' after 'use_device_ptr'}}
+  {}
+#pragma omp target data map(i) use_device_ptr( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
+  {}
+#pragma omp target data map(i) use_device_ptr() // expected-error {{expected expression}}
+  {}
+#pragma omp target data map(i) use_device_ptr(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
+  {}
+#pragma omp target data map(i) use_device_ptr(argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_ptr(i) // expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_ptr(j) // expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_ptr(k) // OK
+  {}
+#pragma omp target data map(i) use_device_ptr(z) // OK
+  {}
+#pragma omp target data map(i) use_device_ptr(aa) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_ptr(e) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_ptr(g) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_ptr(k,i,j) // expected-error2 {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_ptr(d) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_ptr(da) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+  return tmain<int, 3>(argc); // expected-note {{in instantiation of function template specialization 'tmain<int, 3>' requested here}}
+}

tools/libclang/CIndex.cpp

 void OMPClauseEnqueue::VisitOMPFromClause(const OMPFromClause *C) {
   VisitOMPClauseList(C);
 }
+void OMPClauseEnqueue::VisitOMPUseDevicePtrClause(const OMPUseDevicePtrClause *C) {
+  VisitOMPClauseList(C);
+}
 }
 
 void EnqueueVisitor::EnqueueChildren(const OMPClause *S) {