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;
+ /// 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:
+ /// 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);
+ /// 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
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
"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>;
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)
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 {
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);
+}
Node->getDefaultmapKind());
OS << ")";
}
+
+void OMPClausePrinter::VisitOMPUseDevicePtrClause(OMPUseDevicePtrClause *Node) {
+ if (!Node->varlist_empty()) {
+ OS << "use_device_ptr";
+ VisitOMPClauseList(Node, '(');
+ OS << ")";
+ }
+}
}
//===----------------------------------------------------------------------===//
void OMPClauseProfiler::VisitOMPFromClause(const OMPFromClause *C) {
VisitOMPClauseList(C);
}
+void OMPClauseProfiler::VisitOMPUseDevicePtrClause(
+ const OMPUseDevicePtrClause *C) {
+ VisitOMPClauseList(C);
+}
}
void
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");
case OMPC_uniform:
case OMPC_to:
case OMPC_from:
+ case OMPC_use_device_ptr:
llvm_unreachable("Clause is not allowed in 'omp atomic'.");
}
}
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
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);
+}
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
//===----------------------------------------------------------------------===//
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.
//===----------------------------------------------------------------------===//
}
}
+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.
//===----------------------------------------------------------------------===//
--- /dev/null
+// RxUN: %clang_cc1 -verify -fopenmp -std=c++11 -ast-print %s | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+struct ST {
+ int *a;
+};
+struct SA {
+ int i, j;
+ int *k = &j;
+ int *&z = k;
+ void func(int arg) {
+#pragma omp target data map(tofrom: i) use_device_ptr(k)
+ {}
+#pragma omp target data map(tofrom: i) use_device_ptr(z)
+ {}
+ return;
+ }
+};
+// CHECK: struct SA
+// CHECK: void func(
+// CHECK: #pragma omp target data map(tofrom: this->i) use_device_ptr(this->k)
+// CHECK: #pragma omp target data map(tofrom: this->i) use_device_ptr(this->z)
+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>
+T tmain(T argc) {
+ T i;
+ T &j = i;
+ T *k = &j;
+ T *&z = k;
+#pragma omp target data map(tofrom: i) use_device_ptr(k)
+ {}
+#pragma omp target data map(tofrom: i) use_device_ptr(z)
+ {}
+ return 0;
+}
+
+// CHECK: template <typename T = int> int tmain(int argc) {
+// CHECK-NEXT: int i;
+// CHECK-NEXT: int &j = i;
+// CHECK-NEXT: int *k = &j;
+// CHECK-NEXT: int *&z = k;
+// CHECK-NEXT: #pragma omp target data map(tofrom: i) use_device_ptr(k)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target data map(tofrom: i) use_device_ptr(z)
+
+// CHECK: template <typename T = int *> int *tmain(int *argc) {
+// CHECK-NEXT: int *i;
+// CHECK-NEXT: int *&j = i;
+// CHECK-NEXT: int **k = &j;
+// CHECK-NEXT: int **&z = k;
+// CHECK-NEXT: #pragma omp target data map(tofrom: i) use_device_ptr(k)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target data map(tofrom: i) use_device_ptr(z)
+
+// CHECK-LABEL: int main(int argc, char **argv) {
+int main(int argc, char **argv) {
+ int i;
+ int &j = i;
+ int *k = &j;
+ int *&z = k;
+// CHECK-NEXT: int i;
+// CHECK-NEXT: int &j = i;
+// CHECK-NEXT: int *k = &j;
+// CHECK-NEXT: int *&z = k;
+#pragma omp target data map(tofrom: i) use_device_ptr(k)
+// CHECK-NEXT: #pragma omp target data map(tofrom: i) use_device_ptr(k)
+ {}
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+#pragma omp target data map(tofrom: i) use_device_ptr(z)
+// CHECK-NEXT: #pragma omp target data map(tofrom: i) use_device_ptr(z)
+ {}
+ return tmain<int>(argc) + (*tmain<int*>(&argc));
+}
+
+#endif
--- /dev/null
+// 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}}
+}
void OMPClauseEnqueue::VisitOMPFromClause(const OMPFromClause *C) {
VisitOMPClauseList(C);
}
+void OMPClauseEnqueue::VisitOMPUseDevicePtrClause(const OMPUseDevicePtrClause *C) {
+ VisitOMPClauseList(C);
+}
}
void EnqueueVisitor::EnqueueChildren(const OMPClause *S) {
projects / clang / commitdiff