Index: include/clang/AST/DataRecursiveASTVisitor.h =================================================================== --- include/clang/AST/DataRecursiveASTVisitor.h +++ include/clang/AST/DataRecursiveASTVisitor.h @@ -2666,6 +2666,12 @@ return true; } +template +bool RecursiveASTVisitor::VisitOMPDeviceClause(OMPDeviceClause *C) { + TRY_TO(TraverseStmt(C->getDevice())); + 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 Index: include/clang/AST/OpenMPClause.h =================================================================== --- include/clang/AST/OpenMPClause.h +++ include/clang/AST/OpenMPClause.h @@ -2328,6 +2328,63 @@ } }; +/// \brief This represents 'device' clause in the '#pragma omp ...' +/// directive. +/// +/// \code +/// #pragma omp target device(a) +/// \endcode +/// In this example directive '#pragma omp target' has clause 'device' +/// with single expression 'a'. +/// +class OMPDeviceClause : public OMPClause { + friend class OMPClauseReader; + /// \brief Location of '('. + SourceLocation LParenLoc; + /// \brief Device number. + Stmt *Device; + /// \brief Set the device number. + /// + /// \param E Device number. + /// + void setDevice(Expr *E) { Device = E; } + +public: + /// \brief Build 'device' clause. + /// + /// \param E Expression associated with this clause. + /// \param StartLoc Starting location of the clause. + /// \param LParenLoc Location of '('. + /// \param EndLoc Ending location of the clause. + /// + OMPDeviceClause(Expr *E, SourceLocation StartLoc, SourceLocation LParenLoc, + SourceLocation EndLoc) + : OMPClause(OMPC_device, StartLoc, EndLoc), LParenLoc(LParenLoc), + Device(E) {} + + /// \brief Build an empty clause. + /// + OMPDeviceClause() + : OMPClause(OMPC_device, SourceLocation(), SourceLocation()), + LParenLoc(SourceLocation()), Device(0) {} + /// \brief Sets the location of '('. + void setLParenLoc(SourceLocation Loc) { LParenLoc = Loc; } + /// \brief Returns the location of '('. + SourceLocation getLParenLoc() const { return LParenLoc; } + /// \brief Return device number. + Expr *getDevice() { return dyn_cast_or_null(Device); } + /// \brief Return device number. + Expr *getDevice() const { return dyn_cast_or_null(Device); } + + static bool classof(const OMPClause *T) { + return T->getClauseKind() == OMPC_device; + } + + child_range children() { + return child_range(child_iterator(), child_iterator()); + } +}; + } // end namespace clang #endif Index: include/clang/AST/RecursiveASTVisitor.h =================================================================== --- include/clang/AST/RecursiveASTVisitor.h +++ include/clang/AST/RecursiveASTVisitor.h @@ -2698,6 +2698,12 @@ return true; } +template +bool RecursiveASTVisitor::VisitOMPDeviceClause(OMPDeviceClause *C) { + TRY_TO(TraverseStmt(C->getDevice())); + 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 Index: include/clang/Basic/OpenMPKinds.def =================================================================== --- include/clang/Basic/OpenMPKinds.def +++ include/clang/Basic/OpenMPKinds.def @@ -133,6 +133,7 @@ OPENMP_CLAUSE(capture, OMPCaptureClause) OPENMP_CLAUSE(seq_cst, OMPSeqCstClause) OPENMP_CLAUSE(depend, OMPDependClause) +OPENMP_CLAUSE(device, OMPDeviceClause) // Clauses allowed for OpenMP directive 'parallel'. OPENMP_PARALLEL_CLAUSE(if) @@ -275,10 +276,12 @@ // Clauses allowed for OpenMP directive 'target'. // TODO More clauses for 'target' directive. OPENMP_TARGET_CLAUSE(if) +OPENMP_TARGET_CLAUSE(device) // Clauses allowed for OpenMP directive 'target data'. // TODO More clauses for 'target data' directive. OPENMP_TARGET_DATA_CLAUSE(if) +OPENMP_TARGET_DATA_CLAUSE(device) // Clauses allowed for OpenMP directive 'teams'. // TODO More clauses for 'teams' directive. Index: include/clang/Sema/Sema.h =================================================================== --- include/clang/Sema/Sema.h +++ include/clang/Sema/Sema.h @@ -8037,7 +8037,11 @@ SourceLocation ColonLoc, ArrayRef VarList, SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation EndLoc); - + /// \brief Called on well-formed 'device' clause. + OMPClause *ActOnOpenMPDeviceClause(Expr *Device, SourceLocation StartLoc, + SourceLocation LParenLoc, + SourceLocation EndLoc); + /// \brief The kind of conversion being performed. enum CheckedConversionKind { /// \brief An implicit conversion. Index: lib/AST/StmtPrinter.cpp =================================================================== --- lib/AST/StmtPrinter.cpp +++ lib/AST/StmtPrinter.cpp @@ -683,6 +683,12 @@ OS << "seq_cst"; } +void OMPClausePrinter::VisitOMPDeviceClause(OMPDeviceClause *Node) { + OS << "device("; + Node->getDevice()->printPretty(OS, nullptr, Policy, 0); + OS << ")"; +} + template void OMPClausePrinter::VisitOMPClauseList(T *Node, char StartSym) { for (typename T::varlist_iterator I = Node->varlist_begin(), Index: lib/AST/StmtProfile.cpp =================================================================== --- lib/AST/StmtProfile.cpp +++ lib/AST/StmtProfile.cpp @@ -428,6 +428,10 @@ void OMPClauseProfiler::VisitOMPDependClause(const OMPDependClause *C) { VisitOMPClauseList(C); } +void OMPClauseProfiler::VisitOMPDeviceClause(const OMPDeviceClause *C) { + if (C->getDevice()) + Profiler->VisitStmt(C->getDevice()); +} } void Index: lib/Basic/OpenMPKinds.cpp =================================================================== --- lib/Basic/OpenMPKinds.cpp +++ lib/Basic/OpenMPKinds.cpp @@ -122,6 +122,7 @@ case OMPC_update: case OMPC_capture: case OMPC_seq_cst: + case OMPC_device: break; } llvm_unreachable("Invalid OpenMP simple clause kind"); @@ -195,6 +196,7 @@ case OMPC_update: case OMPC_capture: case OMPC_seq_cst: + case OMPC_device: break; } llvm_unreachable("Invalid OpenMP simple clause kind"); Index: lib/CodeGen/CGStmtOpenMP.cpp =================================================================== --- lib/CodeGen/CGStmtOpenMP.cpp +++ lib/CodeGen/CGStmtOpenMP.cpp @@ -2055,6 +2055,7 @@ case OMPC_threadprivate: case OMPC_depend: case OMPC_mergeable: + case OMPC_device: llvm_unreachable("Clause is not allowed in 'omp atomic'."); } } Index: lib/Parse/ParseOpenMP.cpp =================================================================== --- lib/Parse/ParseOpenMP.cpp +++ lib/Parse/ParseOpenMP.cpp @@ -396,7 +396,7 @@ /// lastprivate-clause | reduction-clause | proc_bind-clause | /// schedule-clause | copyin-clause | copyprivate-clause | untied-clause | /// mergeable-clause | flush-clause | read-clause | write-clause | -/// update-clause | capture-clause | seq_cst-clause +/// update-clause | capture-clause | seq_cst-clause | device-clause /// OMPClause *Parser::ParseOpenMPClause(OpenMPDirectiveKind DKind, OpenMPClauseKind CKind, bool FirstClause) { @@ -415,12 +415,15 @@ case OMPC_num_threads: case OMPC_safelen: case OMPC_collapse: + case OMPC_device: // OpenMP [2.5, Restrictions] // At most one if clause can appear on the directive. // At most one num_threads clause can appear on the directive. // OpenMP [2.8.1, simd construct, Restrictions] // Only one safelen clause can appear on a simd directive. // Only one collapse clause can appear on a simd directive. + // OpenMP [2.9.1, target data construct, Restrictions] + // At most one device clause can appear on the directive. // OpenMP [2.11.1, task Construct, Restrictions] // At most one if clause can appear on the directive. // At most one final clause can appear on the directive. Index: lib/Sema/SemaOpenMP.cpp =================================================================== --- lib/Sema/SemaOpenMP.cpp +++ lib/Sema/SemaOpenMP.cpp @@ -4415,6 +4415,9 @@ case OMPC_collapse: Res = ActOnOpenMPCollapseClause(Expr, StartLoc, LParenLoc, EndLoc); break; + case OMPC_device: + Res = ActOnOpenMPDeviceClause(Expr, StartLoc, LParenLoc, EndLoc); + break; case OMPC_default: case OMPC_proc_bind: case OMPC_schedule: @@ -4654,6 +4657,7 @@ case OMPC_capture: case OMPC_seq_cst: case OMPC_depend: + case OMPC_device: case OMPC_unknown: llvm_unreachable("Clause is not allowed."); } @@ -4775,6 +4779,7 @@ case OMPC_capture: case OMPC_seq_cst: case OMPC_depend: + case OMPC_device: case OMPC_unknown: llvm_unreachable("Clause is not allowed."); } @@ -4898,6 +4903,7 @@ case OMPC_threadprivate: case OMPC_flush: case OMPC_depend: + case OMPC_device: case OMPC_unknown: llvm_unreachable("Clause is not allowed."); } @@ -5014,6 +5020,7 @@ case OMPC_update: case OMPC_capture: case OMPC_seq_cst: + case OMPC_device: case OMPC_unknown: llvm_unreachable("Clause is not allowed."); } @@ -6608,3 +6615,28 @@ return OMPDependClause::Create(Context, StartLoc, LParenLoc, EndLoc, DepKind, DepLoc, ColonLoc, Vars); } + +OMPClause *Sema::ActOnOpenMPDeviceClause(Expr *Device, SourceLocation StartLoc, + SourceLocation LParenLoc, + SourceLocation EndLoc) { + Expr *ValExpr = Device; + if (!ValExpr->isTypeDependent() && !ValExpr->isValueDependent() && + !ValExpr->isInstantiationDependent()) { + SourceLocation Loc = ValExpr->getExprLoc(); + ExprResult Value = PerformOpenMPImplicitIntegerConversion(Loc, ValExpr); + if (Value.isInvalid()) + return nullptr; + + // OpenMP [2.9.1, Restrictions] + // The device expression must evaluate to a non-negative integer value. + llvm::APSInt Result; + if (Value.get()->isIntegerConstantExpr(Result, Context) && + Result.isSigned() && !Result.isStrictlyPositive()) { + Diag(Loc, diag::err_omp_negative_expression_in_clause) + << "device" << ValExpr->getSourceRange(); + return nullptr; + } + } + + return new (Context) OMPDeviceClause(ValExpr, StartLoc, LParenLoc, EndLoc); +} Index: lib/Sema/TreeTransform.h =================================================================== --- lib/Sema/TreeTransform.h +++ lib/Sema/TreeTransform.h @@ -1586,6 +1586,17 @@ StartLoc, LParenLoc, EndLoc); } + /// \brief Build a new OpenMP 'device' clause. + /// + /// By default, performs semantic analysis to build the new statement. + /// Subclasses may override this routine to provide different behavior. + OMPClause *RebuildOMPDeviceClause(Expr *Device, SourceLocation StartLoc, + SourceLocation LParenLoc, + SourceLocation EndLoc) { + return getSema().ActOnOpenMPDeviceClause(Device, StartLoc, LParenLoc, + EndLoc); + } + /// \brief Rebuild the operand to an Objective-C \@synchronized statement. /// /// By default, performs semantic analysis to build the new statement. @@ -7480,6 +7491,16 @@ C->getLocStart(), C->getLParenLoc(), C->getLocEnd()); } +template +OMPClause * +TreeTransform::TransformOMPDeviceClause(OMPDeviceClause *C) { + ExprResult E = getDerived().TransformExpr(C->getDevice()); + if (E.isInvalid()) + return nullptr; + return getDerived().RebuildOMPDeviceClause( + E.get(), C->getLocStart(), C->getLParenLoc(), C->getLocEnd()); +} + //===----------------------------------------------------------------------===// // Expression transformation //===----------------------------------------------------------------------===// Index: lib/Serialization/ASTReaderStmt.cpp =================================================================== --- lib/Serialization/ASTReaderStmt.cpp +++ lib/Serialization/ASTReaderStmt.cpp @@ -1788,6 +1788,9 @@ case OMPC_depend: C = OMPDependClause::CreateEmpty(Context, Record[Idx++]); break; + case OMPC_device: + C = new (Context) OMPDeviceClause(); + break; } Visit(C); C->setLocStart(Reader->ReadSourceLocation(Record, Idx)); @@ -2065,6 +2068,11 @@ C->setVarRefs(Vars); } +void OMPClauseReader::VisitOMPDeviceClause(OMPDeviceClause *C) { + C->setDevice(Reader->Reader.ReadSubExpr()); + C->setLParenLoc(Reader->ReadSourceLocation(Record, Idx)); +} + //===----------------------------------------------------------------------===// // OpenMP Directives. //===----------------------------------------------------------------------===// Index: lib/Serialization/ASTWriterStmt.cpp =================================================================== --- lib/Serialization/ASTWriterStmt.cpp +++ lib/Serialization/ASTWriterStmt.cpp @@ -1916,6 +1916,11 @@ Writer->Writer.AddStmt(VE); } +void OMPClauseWriter::VisitOMPDeviceClause(OMPDeviceClause *C) { + Writer->Writer.AddStmt(C->getDevice()); + Writer->Writer.AddSourceLocation(C->getLParenLoc(), Record); +} + //===----------------------------------------------------------------------===// // OpenMP Directives. //===----------------------------------------------------------------------===// Index: test/OpenMP/target_data_device_messages.cpp =================================================================== --- /dev/null +++ test/OpenMP/target_data_device_messages.cpp @@ -0,0 +1,28 @@ +// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp -ferror-limit 100 -o - %s + +void foo() { +} + +bool foobool(int argc) { + return argc; +} + +struct S1; // expected-note {{declared here}} + +int main(int argc, char **argv) { + #pragma omp target data device // expected-error {{expected '(' after 'device'}} + #pragma omp target data device ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target data device () // expected-error {{expected expression}} + #pragma omp target data device (argc // expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target data device (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target data' are ignored}} +#pragma omp target data device (argc > 0 ? argv[1] : argv[2]) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}} + #pragma omp target data device (argc + argc) + #pragma omp target data device (argc), device (argc+1) // expected-error {{directive '#pragma omp target data' cannot contain more than one 'device' clause}} + #pragma omp target data device (S1) // expected-error {{'S1' does not refer to a value}} + #pragma omp target data device (-2) // expected-error {{argument to 'device' clause must be a positive integer value}} + #pragma omp target device (-10u) + #pragma omp target device (3.14) // expected-error {{expression must have integral or unscoped enumeration type, not 'double'}} + foo(); + + return 0; +} Index: test/OpenMP/target_device_messages.cpp =================================================================== --- /dev/null +++ test/OpenMP/target_device_messages.cpp @@ -0,0 +1,28 @@ +// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp -ferror-limit 100 -o - %s + +void foo() { +} + +bool foobool(int argc) { + return argc; +} + +struct S1; // expected-note {{declared here}} + +int main(int argc, char **argv) { + #pragma omp target device // expected-error {{expected '(' after 'device'}} + #pragma omp target device ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target device () // expected-error {{expected expression}} + #pragma omp target device (argc // expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target device (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target' are ignored}} +#pragma omp target device (argc > 0 ? argv[1] : argv[2]) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}} + #pragma omp target device (argc + argc) + #pragma omp target device (argc), device (argc+1) // expected-error {{directive '#pragma omp target' cannot contain more than one 'device' clause}} + #pragma omp target device (S1) // expected-error {{'S1' does not refer to a value}} + #pragma omp target device (-2) // expected-error {{argument to 'device' clause must be a positive integer value}} + #pragma omp target device (-10u) + #pragma omp target device (3.14) // expected-error {{expression must have integral or unscoped enumeration type, not 'double'}} + foo(); + + return 0; +} Index: tools/libclang/CIndex.cpp =================================================================== --- tools/libclang/CIndex.cpp +++ tools/libclang/CIndex.cpp @@ -2051,6 +2051,8 @@ void OMPClauseEnqueue::VisitOMPSeqCstClause(const OMPSeqCstClause *) {} +void OMPClauseEnqueue::VisitOMPDeviceClause(const OMPDeviceClause *C) {} + template void OMPClauseEnqueue::VisitOMPClauseList(T *Node) { for (const auto *I : Node->varlists()) {