[clang] 187ccc6 - [clang][OpenMP5.1] Initial parsing/sema for has_device_addr
Jennifer Yu via cfe-commits
cfe-commits at lists.llvm.org
Fri Apr 8 21:32:18 PDT 2022
Author: Jennifer Yu
Date: 2022-04-08T21:19:38-07:00
New Revision: 187ccc66fa5d0b04189cdbd8266fc386e60f48aa
URL: https://github.com/llvm/llvm-project/commit/187ccc66fa5d0b04189cdbd8266fc386e60f48aa
DIFF: https://github.com/llvm/llvm-project/commit/187ccc66fa5d0b04189cdbd8266fc386e60f48aa.diff
LOG: [clang][OpenMP5.1] Initial parsing/sema for has_device_addr
Added basic parsing/sema/ support for the 'has_device_addr' clause.
Differential Revision: https://reviews.llvm.org/D123402
Added:
clang/test/OpenMP/target_has_device_addr_ast_print.cpp
clang/test/OpenMP/target_has_device_addr_messages.cpp
Modified:
clang/include/clang/AST/OpenMPClause.h
clang/include/clang/AST/RecursiveASTVisitor.h
clang/include/clang/Sema/Sema.h
clang/lib/AST/OpenMPClause.cpp
clang/lib/AST/StmtProfile.cpp
clang/lib/Basic/OpenMPKinds.cpp
clang/lib/CodeGen/CGStmtOpenMP.cpp
clang/lib/Parse/ParseOpenMP.cpp
clang/lib/Sema/SemaOpenMP.cpp
clang/lib/Sema/TreeTransform.h
clang/lib/Serialization/ASTReader.cpp
clang/lib/Serialization/ASTWriter.cpp
clang/test/Analysis/cfg-openmp.cpp
clang/tools/libclang/CIndex.cpp
flang/lib/Semantics/check-omp-structure.cpp
llvm/include/llvm/Frontend/OpenMP/OMP.td
Removed:
################################################################################
diff --git a/clang/include/clang/AST/OpenMPClause.h b/clang/include/clang/AST/OpenMPClause.h
index 3ecc1d40fafc6..3103f61d4248d 100644
--- a/clang/include/clang/AST/OpenMPClause.h
+++ b/clang/include/clang/AST/OpenMPClause.h
@@ -7416,6 +7416,110 @@ class OMPIsDevicePtrClause final
}
};
+/// This represents clause 'has_device_ptr' in the '#pragma omp ...'
+/// directives.
+///
+/// \code
+/// #pragma omp target has_device_addr(a,b)
+/// \endcode
+/// In this example directive '#pragma omp target' has clause
+/// 'has_device_ptr' with the variables 'a' and 'b'.
+class OMPHasDeviceAddrClause final
+ : public OMPMappableExprListClause<OMPHasDeviceAddrClause>,
+ private llvm::TrailingObjects<
+ OMPHasDeviceAddrClause, Expr *, ValueDecl *, unsigned,
+ OMPClauseMappableExprCommon::MappableComponent> {
+ friend class OMPClauseReader;
+ friend OMPMappableExprListClause;
+ friend OMPVarListClause;
+ friend TrailingObjects;
+
+ /// Build clause with number of variables \a NumVars.
+ ///
+ /// \param Locs Locations needed to build a mappable clause. It includes 1)
+ /// StartLoc: starting location of the clause (the clause keyword); 2)
+ /// LParenLoc: location of '('; 3) EndLoc: ending location of the clause.
+ /// \param Sizes All required sizes to build a mappable clause. It includes 1)
+ /// NumVars: number of expressions listed in this clause; 2)
+ /// 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 OMPHasDeviceAddrClause(const OMPVarListLocTy &Locs,
+ const OMPMappableExprListSizeTy &Sizes)
+ : OMPMappableExprListClause(llvm::omp::OMPC_has_device_addr, Locs,
+ Sizes) {}
+
+ /// Build an empty clause.
+ ///
+ /// \param Sizes All required sizes to build a mappable clause. It includes 1)
+ /// NumVars: number of expressions listed in this clause; 2)
+ /// 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 OMPHasDeviceAddrClause(const OMPMappableExprListSizeTy &Sizes)
+ : OMPMappableExprListClause(llvm::omp::OMPC_has_device_addr,
+ OMPVarListLocTy(), Sizes) {}
+
+ /// Define the sizes of each trailing object array except the last one. This
+ /// is required for TrailingObjects to work properly.
+ size_t numTrailingObjects(OverloadToken<Expr *>) const {
+ return varlist_size();
+ }
+ size_t numTrailingObjects(OverloadToken<ValueDecl *>) const {
+ return getUniqueDeclarationsNum();
+ }
+ size_t numTrailingObjects(OverloadToken<unsigned>) const {
+ return getUniqueDeclarationsNum() + getTotalComponentListNum();
+ }
+
+public:
+ /// Creates clause with a list of variables \a Vars.
+ ///
+ /// \param C AST context.
+ /// \param Locs Locations needed to build a mappable clause. It includes 1)
+ /// StartLoc: starting location of the clause (the clause keyword); 2)
+ /// LParenLoc: location of '('; 3) EndLoc: ending location of the clause.
+ /// \param Vars The original expression used in the clause.
+ /// \param Declarations Declarations used in the clause.
+ /// \param ComponentLists Component lists used in the clause.
+ static OMPHasDeviceAddrClause *
+ Create(const ASTContext &C, const OMPVarListLocTy &Locs,
+ ArrayRef<Expr *> Vars, ArrayRef<ValueDecl *> Declarations,
+ MappableExprComponentListsRef ComponentLists);
+
+ /// Creates an empty clause with the place for \a NumVars variables.
+ ///
+ /// \param C AST context.
+ /// \param Sizes All required sizes to build a mappable clause. It includes 1)
+ /// NumVars: number of expressions listed in this clause; 2)
+ /// 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.
+ static OMPHasDeviceAddrClause *
+ CreateEmpty(const ASTContext &C, const OMPMappableExprListSizeTy &Sizes);
+
+ child_range children() {
+ return child_range(reinterpret_cast<Stmt **>(varlist_begin()),
+ reinterpret_cast<Stmt **>(varlist_end()));
+ }
+
+ const_child_range children() const {
+ auto Children = const_cast<OMPHasDeviceAddrClause *>(this)->children();
+ return const_child_range(Children.begin(), Children.end());
+ }
+
+ child_range used_children() {
+ return child_range(child_iterator(), child_iterator());
+ }
+ const_child_range used_children() const {
+ return const_child_range(const_child_iterator(), const_child_iterator());
+ }
+
+ static bool classof(const OMPClause *T) {
+ return T->getClauseKind() == llvm::omp::OMPC_has_device_addr;
+ }
+};
+
/// This represents clause 'nontemporal' in the '#pragma omp ...' directives.
///
/// \code
diff --git a/clang/include/clang/AST/RecursiveASTVisitor.h b/clang/include/clang/AST/RecursiveASTVisitor.h
index 0319bc2e1c988..0fef6982930e4 100644
--- a/clang/include/clang/AST/RecursiveASTVisitor.h
+++ b/clang/include/clang/AST/RecursiveASTVisitor.h
@@ -3702,6 +3702,13 @@ bool RecursiveASTVisitor<Derived>::VisitOMPIsDevicePtrClause(
return true;
}
+template <typename Derived>
+bool RecursiveASTVisitor<Derived>::VisitOMPHasDeviceAddrClause(
+ OMPHasDeviceAddrClause *C) {
+ TRY_TO(VisitOMPClauseList(C));
+ return true;
+}
+
template <typename Derived>
bool RecursiveASTVisitor<Derived>::VisitOMPNontemporalClause(
OMPNontemporalClause *C) {
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index 5ec03391e287b..e7f417a82003b 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -11541,6 +11541,9 @@ class Sema final {
/// Called on well-formed 'is_device_ptr' clause.
OMPClause *ActOnOpenMPIsDevicePtrClause(ArrayRef<Expr *> VarList,
const OMPVarListLocTy &Locs);
+ /// Called on well-formed 'has_device_addr' clause.
+ OMPClause *ActOnOpenMPHasDeviceAddrClause(ArrayRef<Expr *> VarList,
+ const OMPVarListLocTy &Locs);
/// Called on well-formed 'nontemporal' clause.
OMPClause *ActOnOpenMPNontemporalClause(ArrayRef<Expr *> VarList,
SourceLocation StartLoc,
diff --git a/clang/lib/AST/OpenMPClause.cpp b/clang/lib/AST/OpenMPClause.cpp
index 1bd049b880058..9f95c64206028 100644
--- a/clang/lib/AST/OpenMPClause.cpp
+++ b/clang/lib/AST/OpenMPClause.cpp
@@ -146,6 +146,7 @@ const OMPClauseWithPreInit *OMPClauseWithPreInit::get(const OMPClause *C) {
case OMPC_use_device_ptr:
case OMPC_use_device_addr:
case OMPC_is_device_ptr:
+ case OMPC_has_device_addr:
case OMPC_unified_address:
case OMPC_unified_shared_memory:
case OMPC_reverse_offload:
@@ -244,6 +245,7 @@ const OMPClauseWithPostUpdate *OMPClauseWithPostUpdate::get(const OMPClause *C)
case OMPC_use_device_ptr:
case OMPC_use_device_addr:
case OMPC_is_device_ptr:
+ case OMPC_has_device_addr:
case OMPC_unified_address:
case OMPC_unified_shared_memory:
case OMPC_reverse_offload:
@@ -1432,6 +1434,53 @@ OMPIsDevicePtrClause::CreateEmpty(const ASTContext &C,
return new (Mem) OMPIsDevicePtrClause(Sizes);
}
+OMPHasDeviceAddrClause *
+OMPHasDeviceAddrClause::Create(const ASTContext &C, const OMPVarListLocTy &Locs,
+ ArrayRef<Expr *> Vars,
+ ArrayRef<ValueDecl *> Declarations,
+ MappableExprComponentListsRef ComponentLists) {
+ OMPMappableExprListSizeTy Sizes;
+ Sizes.NumVars = Vars.size();
+ Sizes.NumUniqueDeclarations = getUniqueDeclarationsTotalNumber(Declarations);
+ Sizes.NumComponentLists = ComponentLists.size();
+ Sizes.NumComponents = getComponentsTotalNumber(ComponentLists);
+
+ // We need to allocate:
+ // NumVars x Expr* - we have an original list expression for each clause list
+ // entry.
+ // NumUniqueDeclarations x ValueDecl* - unique base declarations associated
+ // with each component list.
+ // (NumUniqueDeclarations + NumComponentLists) x unsigned - we specify the
+ // number of lists for each unique declaration and the size of each component
+ // list.
+ // NumComponents x MappableComponent - the total of all the components in all
+ // the lists.
+ void *Mem = C.Allocate(
+ totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
+ OMPClauseMappableExprCommon::MappableComponent>(
+ Sizes.NumVars, Sizes.NumUniqueDeclarations,
+ Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
+ Sizes.NumComponents));
+
+ auto *Clause = new (Mem) OMPHasDeviceAddrClause(Locs, Sizes);
+
+ Clause->setVarRefs(Vars);
+ Clause->setClauseInfo(Declarations, ComponentLists);
+ return Clause;
+}
+
+OMPHasDeviceAddrClause *
+OMPHasDeviceAddrClause::CreateEmpty(const ASTContext &C,
+ const OMPMappableExprListSizeTy &Sizes) {
+ void *Mem = C.Allocate(
+ totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
+ OMPClauseMappableExprCommon::MappableComponent>(
+ Sizes.NumVars, Sizes.NumUniqueDeclarations,
+ Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
+ Sizes.NumComponents));
+ return new (Mem) OMPHasDeviceAddrClause(Sizes);
+}
+
OMPNontemporalClause *OMPNontemporalClause::Create(const ASTContext &C,
SourceLocation StartLoc,
SourceLocation LParenLoc,
@@ -2259,6 +2308,14 @@ void OMPClausePrinter::VisitOMPIsDevicePtrClause(OMPIsDevicePtrClause *Node) {
}
}
+void OMPClausePrinter::VisitOMPHasDeviceAddrClause(OMPHasDeviceAddrClause *Node) {
+ if (!Node->varlist_empty()) {
+ OS << "has_device_addr";
+ VisitOMPClauseList(Node, '(');
+ OS << ")";
+ }
+}
+
void OMPClausePrinter::VisitOMPNontemporalClause(OMPNontemporalClause *Node) {
if (!Node->varlist_empty()) {
OS << "nontemporal";
diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index 8cf06a2cf992e..77d5d95a1ad10 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -861,6 +861,10 @@ void OMPClauseProfiler::VisitOMPIsDevicePtrClause(
const OMPIsDevicePtrClause *C) {
VisitOMPClauseList(C);
}
+void OMPClauseProfiler::VisitOMPHasDeviceAddrClause(
+ const OMPHasDeviceAddrClause *C) {
+ VisitOMPClauseList(C);
+}
void OMPClauseProfiler::VisitOMPNontemporalClause(
const OMPNontemporalClause *C) {
VisitOMPClauseList(C);
diff --git a/clang/lib/Basic/OpenMPKinds.cpp b/clang/lib/Basic/OpenMPKinds.cpp
index c18d1cc87ac46..2de81dcef3ec0 100644
--- a/clang/lib/Basic/OpenMPKinds.cpp
+++ b/clang/lib/Basic/OpenMPKinds.cpp
@@ -186,6 +186,7 @@ unsigned clang::getOpenMPSimpleClauseType(OpenMPClauseKind Kind, StringRef Str,
case OMPC_use_device_ptr:
case OMPC_use_device_addr:
case OMPC_is_device_ptr:
+ case OMPC_has_device_addr:
case OMPC_unified_address:
case OMPC_unified_shared_memory:
case OMPC_reverse_offload:
@@ -452,6 +453,7 @@ const char *clang::getOpenMPSimpleClauseTypeName(OpenMPClauseKind Kind,
case OMPC_use_device_ptr:
case OMPC_use_device_addr:
case OMPC_is_device_ptr:
+ case OMPC_has_device_addr:
case OMPC_unified_address:
case OMPC_unified_shared_memory:
case OMPC_reverse_offload:
diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp
index a5ee7abc655bd..e85c702907fa0 100644
--- a/clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -6164,6 +6164,7 @@ static void emitOMPAtomicExpr(CodeGenFunction &CGF, OpenMPClauseKind Kind,
case OMPC_use_device_ptr:
case OMPC_use_device_addr:
case OMPC_is_device_ptr:
+ case OMPC_has_device_addr:
case OMPC_unified_address:
case OMPC_unified_shared_memory:
case OMPC_reverse_offload:
diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp
index 60c51d3015987..36105981ec259 100644
--- a/clang/lib/Parse/ParseOpenMP.cpp
+++ b/clang/lib/Parse/ParseOpenMP.cpp
@@ -3108,7 +3108,8 @@ OMPClause *Parser::ParseOpenMPUsesAllocatorClause(OpenMPDirectiveKind DKind) {
/// in_reduction-clause | allocator-clause | allocate-clause |
/// acq_rel-clause | acquire-clause | release-clause | relaxed-clause |
/// depobj-clause | destroy-clause | detach-clause | inclusive-clause |
-/// exclusive-clause | uses_allocators-clause | use_device_addr-clause
+/// exclusive-clause | uses_allocators-clause | use_device_addr-clause |
+/// has_device_addr
///
OMPClause *Parser::ParseOpenMPClause(OpenMPDirectiveKind DKind,
OpenMPClauseKind CKind, bool FirstClause) {
@@ -3290,6 +3291,7 @@ OMPClause *Parser::ParseOpenMPClause(OpenMPDirectiveKind DKind,
case OMPC_use_device_ptr:
case OMPC_use_device_addr:
case OMPC_is_device_ptr:
+ case OMPC_has_device_addr:
case OMPC_allocate:
case OMPC_nontemporal:
case OMPC_inclusive:
@@ -4449,6 +4451,8 @@ bool Parser::ParseOpenMPVarList(OpenMPDirectiveKind DKind,
/// 'use_device_addr' '(' list ')'
/// is_device_ptr-clause:
/// 'is_device_ptr' '(' list ')'
+/// has_device_addr-clause:
+/// 'has_device_addr' '(' list ')'
/// allocate-clause:
/// 'allocate' '(' [ allocator ':' ] list ')'
/// nontemporal-clause:
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index a7b58707c8ecc..9757d0ebb18d8 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -6431,6 +6431,7 @@ StmtResult Sema::ActOnOpenMPExecutableDirective(
case OMPC_use_device_ptr:
case OMPC_use_device_addr:
case OMPC_is_device_ptr:
+ case OMPC_has_device_addr:
case OMPC_nontemporal:
case OMPC_order:
case OMPC_destroy:
@@ -15959,6 +15960,7 @@ OMPClause *Sema::ActOnOpenMPSimpleClause(
case OMPC_use_device_ptr:
case OMPC_use_device_addr:
case OMPC_is_device_ptr:
+ case OMPC_has_device_addr:
case OMPC_unified_address:
case OMPC_unified_shared_memory:
case OMPC_reverse_offload:
@@ -16264,6 +16266,7 @@ OMPClause *Sema::ActOnOpenMPSingleExprWithArgClause(
case OMPC_use_device_ptr:
case OMPC_use_device_addr:
case OMPC_is_device_ptr:
+ case OMPC_has_device_addr:
case OMPC_unified_address:
case OMPC_unified_shared_memory:
case OMPC_reverse_offload:
@@ -16523,6 +16526,7 @@ OMPClause *Sema::ActOnOpenMPClause(OpenMPClauseKind Kind,
case OMPC_use_device_ptr:
case OMPC_use_device_addr:
case OMPC_is_device_ptr:
+ case OMPC_has_device_addr:
case OMPC_atomic_default_mem_order:
case OMPC_device_type:
case OMPC_match:
@@ -17014,6 +17018,9 @@ OMPClause *Sema::ActOnOpenMPVarListClause(
case OMPC_is_device_ptr:
Res = ActOnOpenMPIsDevicePtrClause(VarList, Locs);
break;
+ case OMPC_has_device_addr:
+ Res = ActOnOpenMPHasDeviceAddrClause(VarList, Locs);
+ break;
case OMPC_allocate:
Res = ActOnOpenMPAllocateClause(DepModOrTailExpr, VarList, StartLoc,
LParenLoc, ColonLoc, EndLoc);
@@ -22426,6 +22433,88 @@ OMPClause *Sema::ActOnOpenMPIsDevicePtrClause(ArrayRef<Expr *> VarList,
MVLI.VarComponents);
}
+OMPClause *Sema::ActOnOpenMPHasDeviceAddrClause(ArrayRef<Expr *> VarList,
+ const OMPVarListLocTy &Locs) {
+ MappableVarListInfo MVLI(VarList);
+ for (Expr *RefExpr : VarList) {
+ assert(RefExpr && "NULL expr in OpenMP has_device_addr clause.");
+ SourceLocation ELoc;
+ SourceRange ERange;
+ Expr *SimpleRefExpr = RefExpr;
+ auto Res = getPrivateItem(*this, SimpleRefExpr, ELoc, ERange,
+ /*AllowArraySection=*/true);
+ if (Res.second) {
+ // It will be analyzed later.
+ MVLI.ProcessedVarList.push_back(RefExpr);
+ }
+ ValueDecl *D = Res.first;
+ if (!D)
+ continue;
+
+ // Check if the declaration in the clause does not show up in any data
+ // sharing attribute.
+ DSAStackTy::DSAVarData DVar = DSAStack->getTopDSA(D, /*FromParent=*/false);
+ if (isOpenMPPrivate(DVar.CKind)) {
+ Diag(ELoc, diag::err_omp_variable_in_given_clause_and_dsa)
+ << getOpenMPClauseName(DVar.CKind)
+ << getOpenMPClauseName(OMPC_has_device_addr)
+ << getOpenMPDirectiveName(DSAStack->getCurrentDirective());
+ reportOriginalDsa(*this, DSAStack, D, DVar);
+ continue;
+ }
+
+ const Expr *ConflictExpr;
+ if (DSAStack->checkMappableExprComponentListsForDecl(
+ D, /*CurrentRegionOnly=*/true,
+ [&ConflictExpr](
+ OMPClauseMappableExprCommon::MappableExprComponentListRef R,
+ OpenMPClauseKind) -> bool {
+ ConflictExpr = R.front().getAssociatedExpression();
+ return true;
+ })) {
+ Diag(ELoc, diag::err_omp_map_shared_storage) << RefExpr->getSourceRange();
+ Diag(ConflictExpr->getExprLoc(), diag::note_used_here)
+ << ConflictExpr->getSourceRange();
+ continue;
+ }
+
+ // Store the components in the stack so that they can be used to check
+ // against other clauses later on.
+ OMPClauseMappableExprCommon::MappableComponent MC(
+ SimpleRefExpr, D, /*IsNonContiguous=*/false);
+ DSAStack->addMappableExpressionComponents(
+ D, MC, /*WhereFoundClauseKind=*/OMPC_has_device_addr);
+
+ // Record the expression we've just processed.
+ auto *VD = dyn_cast<VarDecl>(D);
+ if (!VD && !CurContext->isDependentContext()) {
+ DeclRefExpr *Ref =
+ buildCapture(*this, D, SimpleRefExpr, /*WithInit=*/true);
+ assert(Ref && "has_device_addr capture failed");
+ MVLI.ProcessedVarList.push_back(Ref);
+ } else
+ MVLI.ProcessedVarList.push_back(RefExpr->IgnoreParens());
+
+ // Create a mappable component for the list item. List items in this clause
+ // only need a component. We use a null declaration to signal fields in
+ // 'this'.
+ assert((isa<DeclRefExpr>(SimpleRefExpr) ||
+ isa<CXXThisExpr>(cast<MemberExpr>(SimpleRefExpr)->getBase())) &&
+ "Unexpected device pointer expression!");
+ MVLI.VarBaseDeclarations.push_back(
+ isa<DeclRefExpr>(SimpleRefExpr) ? D : nullptr);
+ MVLI.VarComponents.resize(MVLI.VarComponents.size() + 1);
+ MVLI.VarComponents.back().push_back(MC);
+ }
+
+ if (MVLI.ProcessedVarList.empty())
+ return nullptr;
+
+ return OMPHasDeviceAddrClause::Create(Context, Locs, MVLI.ProcessedVarList,
+ MVLI.VarBaseDeclarations,
+ MVLI.VarComponents);
+}
+
OMPClause *Sema::ActOnOpenMPAllocateClause(
Expr *Allocator, ArrayRef<Expr *> VarList, SourceLocation StartLoc,
SourceLocation ColonLoc, SourceLocation LParenLoc, SourceLocation EndLoc) {
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index 773cc668f3378..5eff7cfd7a253 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -2101,6 +2101,15 @@ class TreeTransform {
return getSema().ActOnOpenMPIsDevicePtrClause(VarList, Locs);
}
+ /// Build a new OpenMP 'has_device_addr' clause.
+ ///
+ /// By default, performs semantic analysis to build the new OpenMP clause.
+ /// Subclasses may override this routine to provide
diff erent behavior.
+ OMPClause *RebuildOMPHasDeviceAddrClause(ArrayRef<Expr *> VarList,
+ const OMPVarListLocTy &Locs) {
+ return getSema().ActOnOpenMPHasDeviceAddrClause(VarList, Locs);
+ }
+
/// Build a new OpenMP 'defaultmap' clause.
///
/// By default, performs semantic analysis to build the new OpenMP clause.
@@ -10270,6 +10279,21 @@ TreeTransform<Derived>::TransformOMPIsDevicePtrClause(OMPIsDevicePtrClause *C) {
return getDerived().RebuildOMPIsDevicePtrClause(Vars, Locs);
}
+template <typename Derived>
+OMPClause *TreeTransform<Derived>::TransformOMPHasDeviceAddrClause(
+ OMPHasDeviceAddrClause *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());
+ }
+ OMPVarListLocTy Locs(C->getBeginLoc(), C->getLParenLoc(), C->getEndLoc());
+ return getDerived().RebuildOMPHasDeviceAddrClause(Vars, Locs);
+}
+
template <typename Derived>
OMPClause *
TreeTransform<Derived>::TransformOMPNontemporalClause(OMPNontemporalClause *C) {
diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp
index 33e59fb732828..3a32027023328 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -11867,6 +11867,15 @@ OMPClause *OMPClauseReader::readClause() {
C = OMPIsDevicePtrClause::CreateEmpty(Context, Sizes);
break;
}
+ case llvm::omp::OMPC_has_device_addr: {
+ OMPMappableExprListSizeTy Sizes;
+ Sizes.NumVars = Record.readInt();
+ Sizes.NumUniqueDeclarations = Record.readInt();
+ Sizes.NumComponentLists = Record.readInt();
+ Sizes.NumComponents = Record.readInt();
+ C = OMPHasDeviceAddrClause::CreateEmpty(Context, Sizes);
+ break;
+ }
case llvm::omp::OMPC_allocate:
C = OMPAllocateClause::CreateEmpty(Context, Record.readInt());
break;
@@ -12827,6 +12836,49 @@ void OMPClauseReader::VisitOMPIsDevicePtrClause(OMPIsDevicePtrClause *C) {
C->setComponents(Components, ListSizes);
}
+void OMPClauseReader::VisitOMPHasDeviceAddrClause(OMPHasDeviceAddrClause *C) {
+ C->setLParenLoc(Record.readSourceLocation());
+ auto NumVars = C->varlist_size();
+ auto UniqueDecls = C->getUniqueDeclarationsNum();
+ auto TotalLists = C->getTotalComponentListNum();
+ auto TotalComponents = C->getTotalComponentsNum();
+
+ SmallVector<Expr *, 16> Vars;
+ Vars.reserve(NumVars);
+ for (unsigned I = 0; I != NumVars; ++I)
+ Vars.push_back(Record.readSubExpr());
+ C->setVarRefs(Vars);
+ Vars.clear();
+
+ SmallVector<ValueDecl *, 16> Decls;
+ Decls.reserve(UniqueDecls);
+ for (unsigned I = 0; I < UniqueDecls; ++I)
+ Decls.push_back(Record.readDeclAs<ValueDecl>());
+ C->setUniqueDecls(Decls);
+
+ SmallVector<unsigned, 16> ListsPerDecl;
+ ListsPerDecl.reserve(UniqueDecls);
+ for (unsigned I = 0; I < UniqueDecls; ++I)
+ ListsPerDecl.push_back(Record.readInt());
+ C->setDeclNumLists(ListsPerDecl);
+
+ SmallVector<unsigned, 32> ListSizes;
+ ListSizes.reserve(TotalLists);
+ for (unsigned i = 0; i < TotalLists; ++i)
+ ListSizes.push_back(Record.readInt());
+ C->setComponentListSizes(ListSizes);
+
+ SmallVector<OMPClauseMappableExprCommon::MappableComponent, 32> Components;
+ Components.reserve(TotalComponents);
+ for (unsigned I = 0; I < TotalComponents; ++I) {
+ Expr *AssociatedExpr = Record.readSubExpr();
+ auto *AssociatedDecl = Record.readDeclAs<ValueDecl>();
+ Components.emplace_back(AssociatedExpr, AssociatedDecl,
+ /*IsNonContiguous=*/false);
+ }
+ C->setComponents(Components, ListSizes);
+}
+
void OMPClauseReader::VisitOMPNontemporalClause(OMPNontemporalClause *C) {
C->setLParenLoc(Record.readSourceLocation());
unsigned NumVars = C->varlist_size();
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index 5fd4e8fb27cf8..4e8542e13ff70 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -6831,6 +6831,26 @@ void OMPClauseWriter::VisitOMPIsDevicePtrClause(OMPIsDevicePtrClause *C) {
}
}
+void OMPClauseWriter::VisitOMPHasDeviceAddrClause(OMPHasDeviceAddrClause *C) {
+ Record.push_back(C->varlist_size());
+ Record.push_back(C->getUniqueDeclarationsNum());
+ Record.push_back(C->getTotalComponentListNum());
+ Record.push_back(C->getTotalComponentsNum());
+ Record.AddSourceLocation(C->getLParenLoc());
+ for (auto *E : C->varlists())
+ Record.AddStmt(E);
+ for (auto *D : C->all_decls())
+ Record.AddDeclRef(D);
+ for (auto N : C->all_num_lists())
+ Record.push_back(N);
+ for (auto N : C->all_lists_sizes())
+ Record.push_back(N);
+ for (auto &M : C->all_components()) {
+ Record.AddStmt(M.getAssociatedExpression());
+ Record.AddDeclRef(M.getAssociatedDeclaration());
+ }
+}
+
void OMPClauseWriter::VisitOMPUnifiedAddressClause(OMPUnifiedAddressClause *) {}
void OMPClauseWriter::VisitOMPUnifiedSharedMemoryClause(
diff --git a/clang/test/Analysis/cfg-openmp.cpp b/clang/test/Analysis/cfg-openmp.cpp
index 11325deb20aca..4a6923a855ca0 100644
--- a/clang/test/Analysis/cfg-openmp.cpp
+++ b/clang/test/Analysis/cfg-openmp.cpp
@@ -1,5 +1,126 @@
// RUN: %clang_analyze_cc1 -analyzer-checker=debug.DumpCFG %s 2>&1 -fopenmp -fopenmp-version=45 | FileCheck %s
+// RUN: %clang_analyze_cc1 -analyzer-checker=debug.DumpCFG %s 2>&1 -fopenmp -fopenmp-version=51 | FileCheck %s --check-prefix=OMP51
+
+#if _OPENMP == 202011
+
+// OMP51-LABEL: void target_has_device_addr(int argc)
+void target_has_device_addr(int argc) {
+// OMP51: [B1]
+// OMP51-NEXT: [[#TTD:]]: 5
+// OMP51-NEXT: [[#TTD+1]]: int x = 5;
+// OMP51-NEXT: [[#TTD+2]]: x
+// OMP51-NEXT: [[#TTD+3]]: [B1.[[#TTD+2]]] (ImplicitCastExpr, LValueToRValue, int)
+// OMP51-NEXT: [[#TTD+4]]: [B1.[[#TTD+6]]]
+// OMP51-NEXT: [[#TTD+5]]: [B1.[[#TTD+6]]] = [B1.[[#TTD+3]]]
+// OMP51-NEXT: [[#TTD+6]]: argc
+// OMP51-NEXT: [[#TTD+7]]: #pragma omp target has_device_addr(x)
+// OMP51-NEXT: [B1.[[#TTD+5]]]
+ int x = 5;
+#pragma omp target has_device_addr(x)
+ argc = x;
+}
+// OMP51-LABEL: void target_s_has_device_addr(int argc)
+void target_s_has_device_addr(int argc) {
+ int x, cond, fp, rd, lin, step, map;
+// OMP51-DAG: [B3]
+// OMP51-DAG: [[#TSB:]]: x
+// OMP51-DAG: [[#TSB+1]]: [B3.[[#TSB]]] (ImplicitCastExpr, LValueToRValue, int)
+// OMP51-DAG: [[#TSB+2]]: argc
+// OMP51-DAG: [[#TSB+3]]: [B3.[[#TSB+2]]] = [B3.[[#TSB+1]]]
+// OMP51-DAG: [B1]
+// OMP51-DAG: [[#TS:]]: cond
+// OMP51-DAG: [[#TS+1]]: [B1.[[#TS]]] (ImplicitCastExpr, LValueToRValue, int)
+// OMP51-DAG: [[#TS+2]]: [B1.[[#TS+1]]] (ImplicitCastExpr, IntegralToBoolean, _Bool)
+// OMP51-DAG: [[#TS+3]]: fp
+// OMP51-DAG: [[#TS+4]]: rd
+// OMP51-DAG: [[#TS+5]]: lin
+// OMP51-DAG: [[#TS+6]]: step
+// OMP51-DAG: [[#TS+7]]: [B1.[[#TS+6]]] (ImplicitCastExpr, LValueToRValue, int)
+// OMP51-DAG: [[#TS+8]]: [B3.[[#TSB+2]]]
+// OMP51-DAG: [[#TS+9]]: [B3.[[#TSB]]]
+// OMP51-DAG: [[#TS+10]]: #pragma omp target simd if(cond) firstprivate(fp) reduction(+: rd) linear(lin: step) has_device_addr(map)
+// OMP51-DAG: for (int i = 0;
+// OMP51-DAG: [B3.[[#TSB+3]]];
+#pragma omp target simd if(cond) firstprivate(fp) reduction(+:rd) linear(lin: step) has_device_addr(map)
+ for (int i = 0; i < 10; ++i)
+ argc = x;
+}
+// OMP51-LABEL: void target_t_l_has_device_addr(int argc)
+void target_t_l_has_device_addr(int argc) {
+int x, cond, fp, rd, map;
+// OMP51-DAG: [B3]
+// OMP51-DAG: [[#TTDB:]]: x
+// OMP51-DAG: [[#TTDB+1]]: [B3.[[#TTDB]]] (ImplicitCastExpr, LValueToRValue, int)
+// OMP51-DAG: [[#TTDB+2]]: argc
+// OMP51-DAG: [[#TTDB+3]]: [B3.[[#TTDB+2]]] = [B3.[[#TTDB+1]]]
+// OMP51-DAG: [B1]
+// OMP51-DAG: [[#TTD:]]: cond
+// OMP51-DAG: [[#TTD+1]]: [B1.[[#TTD]]] (ImplicitCastExpr, LValueToRValue, int)
+// OMP51-DAG: [[#TTD+2]]: [B1.[[#TTD+1]]] (ImplicitCastExpr, IntegralToBoolean, _Bool)
+// OMP51-DAG: [[#TTD+3]]: fp
+// OMP51-DAG: [[#TTD+4]]: rd
+// OMP51-DAG: [[#TTD+5]]: [B3.[[#TTDB+2]]]
+// OMP51-DAG: [[#TTD+6]]: [B3.[[#TTDB]]]
+// OMP51-DAG: [[#TTD+7]]: #pragma omp target teams loop if(cond) firstprivate(fp) reduction(+: rd) has_device_addr(map)
+// OMP51-DAG: for (int i = 0;
+// OMP51-DAG: [B3.[[#TTDB+3]]];
+#pragma omp target teams loop if(cond) firstprivate(fp) reduction(+:rd) has_device_addr(map)
+ for (int i = 0; i <10; ++i)
+ argc = x;
+}
+// OMP51-LABEL: void target_p_l_has_device_addr(int argc)
+void target_p_l_has_device_addr(int argc) {
+int x, cond, fp, rd, map;
+#pragma omp target parallel loop if(cond) firstprivate(fp) reduction(+:rd) has_device_addr(map)
+// OMP51-DAG: [B3]
+// OMP51-DAG: [[#TTDB:]]: x
+// OMP51-DAG: [[#TTDB+1]]: [B3.[[#TTDB]]] (ImplicitCastExpr, LValueToRValue, int)
+// OMP51-DAG: [[#TTDB+2]]: argc
+// OMP51-DAG: [[#TTDB+3]]: [B3.[[#TTDB+2]]] = [B3.[[#TTDB+1]]]
+// OMP51-DAG: [B1]
+// OMP51-DAG: [[#TTD:]]: cond
+// OMP51-DAG: [[#TTD+1]]: [B1.[[#TTD]]] (ImplicitCastExpr, LValueToRValue, int)
+// OMP51-DAG: [[#TTD+2]]: [B1.[[#TTD+1]]] (ImplicitCastExpr, IntegralToBoolean, _Bool)
+// OMP51-DAG: [[#TTD+3]]: fp
+// OMP51-DAG: [[#TTD+4]]: rd
+// OMP51-DAG: [[#TTD+5]]: [B3.[[#TTDB+2]]]
+// OMP51-DAG: [[#TTD+6]]: [B3.[[#TTDB]]]
+// OMP51-DAG: [[#TTD+7]]: #pragma omp target parallel loop if(cond) firstprivate(fp) reduction(+: rd) has_device_addr(map)
+// OMP51-DAG: for (int i = 0;
+// OMP51-DAG: [B3.[[#TTDB+3]]];
+ for (int i = 0; i < 10; ++i)
+ argc = x;
+}
+struct SomeKernel {
+ int targetDev;
+ float devPtr;
+ SomeKernel();
+ ~SomeKernel();
+// OMP51-LABEL: template<> void apply<32U>()
+ template<unsigned int nRHS>
+ void apply() {
+// OMP51-DAG: [B1]
+// OMP51-DAG: [[#TTD:]]: 10
+// OMP51-DAG: [[#TTD+1]]: [B1.[[#TTD:]]] (ImplicitCastExpr, IntegralToFloating, float)
+// OMP51-DAG: [[#TTD+2]]: this
+// OMP51-DAG: [[#TTD+3]]: [B1.[[#TTD+2]]]->devPtr
+// OMP51-DAG: [[#TTD+4]]: [B1.[[#TTD+3]]] = [B1.[[#TTD+1]]]
+// OMP51-DAG: [[#TTD+5]]: #pragma omp target has_device_addr(this->devPtr) device(this->targetDev)
+// OMP51-DAG: {
+// OMP51-DAG: [B1.[[#TTD+4]]];
+ #pragma omp target has_device_addr(devPtr) device(targetDev)
+ {
+ devPtr = 10;
+ }
+ }
+};
+void use_template() {
+ SomeKernel aKern;
+ aKern.apply<32>();
+}
+#else // _OPENMP
+
// CHECK-LABEL: void xxx(int argc)
void xxx(int argc) {
// CHECK: [B1]
@@ -771,3 +892,5 @@ void targetparallelloop(int argc) {
for (int i = 0; i < 10; ++i)
argc = x;
}
+
+#endif // _OPENMP
diff --git a/clang/test/OpenMP/target_has_device_addr_ast_print.cpp b/clang/test/OpenMP/target_has_device_addr_ast_print.cpp
new file mode 100644
index 0000000000000..fc8017e951a2f
--- /dev/null
+++ b/clang/test/OpenMP/target_has_device_addr_ast_print.cpp
@@ -0,0 +1,338 @@
+// RUN: %clang_cc1 -verify -fopenmp -std=c++11 -fopenmp-version=51 \
+// RUN: -ast-print %s | FileCheck %s
+
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -fopenmp-version=51 \
+// RUN: -emit-pch -o %t %s
+
+// RUN: %clang_cc1 -fopenmp -std=c++11 -fopenmp-version=51 \
+// RUN: -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+
+// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=51 \
+// RUN: -std=c++11 -ast-print %s | FileCheck %s
+
+// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 \
+// RUN: -fopenmp-version=51 -emit-pch -o %t %s
+
+// RUN: %clang_cc1 -fopenmp-simd -std=c++11 -fopenmp-version=51 \
+// RUN: -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+struct ST {
+ int *a;
+};
+typedef int arr[10];
+typedef ST STarr[10];
+struct SA {
+ const int da[5] = { 0 };
+ ST g[10];
+ STarr &rg = g;
+ int i;
+ int &j = i;
+ int *k = &j;
+ int *&z = k;
+ int aa[10];
+ arr &raa = aa;
+ void func(int arg) {
+#pragma omp target has_device_addr(k)
+ {}
+#pragma omp target has_device_addr(z)
+ {}
+#pragma omp target has_device_addr(aa) // OK
+ {}
+#pragma omp target has_device_addr(raa) // OK
+ {}
+#pragma omp target has_device_addr(g) // OK
+ {}
+#pragma omp target has_device_addr(rg) // OK
+ {}
+#pragma omp target has_device_addr(da) // OK
+ {}
+ return;
+ }
+};
+// CHECK: struct SA
+// CHECK-NEXT: const int da[5] = {0};
+// CHECK-NEXT: ST g[10];
+// CHECK-NEXT: STarr &rg = this->g;
+// CHECK-NEXT: int i;
+// CHECK-NEXT: int &j = this->i;
+// CHECK-NEXT: int *k = &this->j;
+// CHECK-NEXT: int *&z = this->k;
+// CHECK-NEXT: int aa[10];
+// CHECK-NEXT: arr &raa = this->aa;
+// CHECK-NEXT: func(
+// CHECK-NEXT: #pragma omp target has_device_addr(this->k)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(this->z)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(this->aa)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(this->raa)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(this->g)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(this->rg)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(this->da)
+
+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 struct {
+ int a;
+} S6;
+
+template <typename T>
+T tmain(T argc) {
+ const T da[5] = { 0 };
+ S6 h[10];
+ auto &rh = h;
+ T i;
+ T &j = i;
+ T *k = &j;
+ T *&z = k;
+ T aa[10];
+ auto &raa = aa;
+#pragma omp target has_device_addr(k)
+ {}
+#pragma omp target has_device_addr(z)
+ {}
+#pragma omp target has_device_addr(aa)
+ {}
+#pragma omp target has_device_addr(raa)
+ {}
+#pragma omp target has_device_addr(h)
+ {}
+#pragma omp target has_device_addr(rh)
+ {}
+#pragma omp target has_device_addr(da)
+ {}
+ return 0;
+}
+
+// CHECK: template<> int tmain<int>(int argc) {
+// CHECK-NEXT: const int da[5] = {0};
+// CHECK-NEXT: S6 h[10];
+// CHECK-NEXT: auto &rh = h;
+// CHECK-NEXT: int i;
+// CHECK-NEXT: int &j = i;
+// CHECK-NEXT: int *k = &j;
+// CHECK-NEXT: int *&z = k;
+// CHECK-NEXT: int aa[10];
+// CHECK-NEXT: auto &raa = aa;
+// CHECK-NEXT: #pragma omp target has_device_addr(k)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(z)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(aa)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(raa)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(h)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(rh)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(da)
+
+// CHECK: template<> int *tmain<int *>(int *argc) {
+// CHECK-NEXT: int *const da[5] = {0};
+// CHECK-NEXT: S6 h[10];
+// CHECK-NEXT: auto &rh = h;
+// CHECK-NEXT: int *i;
+// CHECK-NEXT: int *&j = i;
+// CHECK-NEXT: int **k = &j;
+// CHECK-NEXT: int **&z = k;
+// CHECK-NEXT: int *aa[10];
+// CHECK-NEXT: auto &raa = aa;
+// CHECK-NEXT: #pragma omp target has_device_addr(k)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(z)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(aa)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(raa)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(h)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(rh)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(da)
+
+// CHECK-LABEL: int main(int argc, char **argv) {
+int main(int argc, char **argv) {
+ const int da[5] = { 0 };
+ S6 h[10];
+ auto &rh = h;
+ int i;
+ int &j = i;
+ int *k = &j;
+ int *&z = k;
+ int aa[10];
+ auto &raa = aa;
+// CHECK-NEXT: const int da[5] = {0};
+// CHECK-NEXT: S6 h[10];
+// CHECK-NEXT: auto &rh = h;
+// CHECK-NEXT: int i;
+// CHECK-NEXT: int &j = i;
+// CHECK-NEXT: int *k = &j;
+// CHECK-NEXT: int *&z = k;
+// CHECK-NEXT: int aa[10];
+// CHECK-NEXT: auto &raa = aa;
+#pragma omp target has_device_addr(k)
+// CHECK-NEXT: #pragma omp target has_device_addr(k)
+ {}
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+#pragma omp target has_device_addr(z)
+// CHECK-NEXT: #pragma omp target has_device_addr(z)
+ {}
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+#pragma omp target has_device_addr(aa)
+// CHECK-NEXT: #pragma omp target has_device_addr(aa)
+ {}
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+#pragma omp target has_device_addr(raa)
+// CHECK-NEXT: #pragma omp target has_device_addr(raa)
+ {}
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+#pragma omp target has_device_addr(h)
+// CHECK-NEXT: #pragma omp target has_device_addr(h)
+ {}
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+#pragma omp target has_device_addr(rh)
+// CHECK-NEXT: #pragma omp target has_device_addr(rh)
+ {}
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+#pragma omp target has_device_addr(da)
+// CHECK-NEXT: #pragma omp target has_device_addr(da)
+ {}
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(da[1:3])
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+#pragma omp target has_device_addr(da[1:3])
+ {}
+ return tmain<int>(argc) + *tmain<int *>(&argc);
+}
+
+struct SomeKernel {
+ int targetDev;
+ float devPtr;
+ SomeKernel();
+ ~SomeKernel();
+
+ template<unsigned int nRHS>
+ void apply() {
+ #pragma omp target has_device_addr(devPtr) device(targetDev)
+ {
+ }
+// CHECK: #pragma omp target has_device_addr(this->devPtr) device(this->targetDev)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+ }
+// CHECK: template<> void apply<32U>() {
+// CHECK: #pragma omp target has_device_addr(this->devPtr) device(this->targetDev)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+};
+
+void use_template() {
+ SomeKernel aKern;
+ aKern.apply<32>();
+}
+#endif
diff --git a/clang/test/OpenMP/target_has_device_addr_messages.cpp b/clang/test/OpenMP/target_has_device_addr_messages.cpp
new file mode 100644
index 0000000000000..52c849f66bc84
--- /dev/null
+++ b/clang/test/OpenMP/target_has_device_addr_messages.cpp
@@ -0,0 +1,273 @@
+// RUN: %clang_cc1 -std=c++11 -fopenmp-version=51 -verify \
+// RUN: -fopenmp -ferror-limit 200 %s -Wuninitialized
+
+// RUN: %clang_cc1 -std=c++11 -fopenmp-version=51 -verify \
+// RUN: -fopenmp-simd -ferror-limit 200 %s -Wuninitialized
+
+struct ST {
+ int *a;
+};
+typedef int arr[10];
+typedef ST STarr[10];
+struct SA {
+ const int d = 5;
+ const int da[5] = { 0 };
+ ST e;
+ ST g[10];
+ STarr &rg = g;
+ int i;
+ int &j = i;
+ int *k = &j;
+ int *&z = k;
+ int aa[10];
+ arr &raa = aa;
+ void func(int arg) {
+#pragma omp target has_device_addr // expected-error {{expected '(' after 'has_device_addr'}}
+ {}
+#pragma omp target has_device_addr( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
+ {}
+#pragma omp target has_device_addr() // expected-error {{expected expression}}
+ {}
+#pragma omp target has_device_addr(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
+ {}
+#pragma omp target has_device_addr(arg // expected-error {{expected ')'}} expected-note {{to match this '('}}
+ {}
+#pragma omp target has_device_addr(i) // OK
+ {}
+#pragma omp target has_device_addr(j) // OK
+ {}
+#pragma omp target has_device_addr(k) // OK
+ {}
+#pragma omp target has_device_addr(z) // OK
+ {}
+#pragma omp target has_device_addr(aa) // OK
+ {}
+#pragma omp target has_device_addr(raa) // OK
+ {}
+#pragma omp target has_device_addr(e) // OK
+ {}
+#pragma omp target has_device_addr(g) // OK
+ {}
+#pragma omp target has_device_addr(rg) // OK
+ {}
+#pragma omp target has_device_addr(k,i,j) // OK
+ {}
+#pragma omp target has_device_addr(d) // OK
+ {}
+#pragma omp target has_device_addr(da) // OK
+ {}
+ 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 struct {
+ int a;
+} S6;
+
+template <typename T, int I>
+T tmain(T argc) {
+ const T d = 5;
+ const T da[5] = { 0 };
+ S4 e(4);
+ S5 g(5);
+ S6 h[10];
+ auto &rh = h;
+ T i;
+ T &j = i;
+ T *k = &j;
+ T *&z = k;
+ T aa[10];
+ auto &raa = aa;
+ S6 *ps;
+#pragma omp target has_device_addr // expected-error {{expected '(' after 'has_device_addr'}}
+ {}
+#pragma omp target has_device_addr( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
+ {}
+#pragma omp target has_device_addr() // expected-error {{expected expression}}
+ {}
+#pragma omp target has_device_addr(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
+ {}
+#pragma omp target has_device_addr(argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+ {}
+#pragma omp target has_device_addr(i) // OK
+ {}
+#pragma omp target has_device_addr(j) // OK
+ {}
+#pragma omp target has_device_addr(k) // OK
+ {}
+#pragma omp target has_device_addr(z) // OK
+ {}
+#pragma omp target has_device_addr(aa) // OK
+ {}
+#pragma omp target has_device_addr(raa) // OK
+ {}
+#pragma omp target has_device_addr(e) // OK
+ {}
+#pragma omp target has_device_addr(g) // OK
+ {}
+#pragma omp target has_device_addr(h) // OK
+ {}
+#pragma omp target has_device_addr(rh) // OK
+ {}
+#pragma omp target has_device_addr(k,i,j) // OK
+ {}
+#pragma omp target has_device_addr(d) // OK
+ {}
+#pragma omp target has_device_addr(da) // OK
+ {}
+#pragma omp target map(ps) has_device_addr(ps) // expected-error{{variable already marked as mapped in current construct}} expected-note{{used here}}
+ {}
+#pragma omp target has_device_addr(ps) map(ps) // expected-error{{variable already marked as mapped in current construct}} expected-note{{used here}}
+ {}
+#pragma omp target map(ps->a) has_device_addr(ps) // expected-error{{variable already marked as mapped in current construct}} expected-note{{used here}}
+ {}
+#pragma omp target has_device_addr(ps) map(ps->a) // expected-error{{pointer cannot be mapped along with a section derived from itself}} expected-note{{used here}}
+ {}
+#pragma omp target has_device_addr(ps) firstprivate(ps) // expected-error{{firstprivate variable cannot be in a has_device_addr clause in '#pragma omp target' directive}}
+ {}
+#pragma omp target firstprivate(ps) has_device_addr(ps) // expected-error{{firstprivate variable cannot be in a has_device_addr clause in '#pragma omp target' directive}} expected-note{{defined as firstprivate}}
+ {}
+#pragma omp target has_device_addr(ps) private(ps) // expected-error{{private variable cannot be in a has_device_addr clause in '#pragma omp target' directive}}
+ {}
+#pragma omp target private(ps) has_device_addr(ps) // expected-error{{private variable cannot be in a has_device_addr clause in '#pragma omp target' directive}} expected-note{{defined as private}}
+ {}
+ return 0;
+}
+
+int main(int argc, char **argv) {
+ const int d = 5;
+ const int da[5] = { 0 };
+ S4 e(4);
+ S5 g(5);
+ S6 h[10];
+ auto &rh = h;
+ int i;
+ int &j = i;
+ int *k = &j;
+ int *&z = k;
+ int aa[10];
+ auto &raa = aa;
+ S6 *ps;
+#pragma omp target has_device_addr // expected-error {{expected '(' after 'has_device_addr'}}
+ {}
+#pragma omp target has_device_addr( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
+ {}
+#pragma omp target has_device_addr() // expected-error {{expected expression}}
+ {}
+#pragma omp target has_device_addr(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
+ {}
+#pragma omp target has_device_addr(argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+ {}
+#pragma omp target has_device_addr(i) // OK
+ {}
+#pragma omp target has_device_addr(j) // OK
+ {}
+#pragma omp target has_device_addr(k) // OK
+ {}
+#pragma omp target has_device_addr(z) // OK
+ {}
+#pragma omp target has_device_addr(aa) // OK
+ {}
+#pragma omp target has_device_addr(raa) // OK
+ {}
+#pragma omp target has_device_addr(e) // OK
+ {}
+#pragma omp target has_device_addr(g) // OK
+ {}
+#pragma omp target has_device_addr(h) // OK
+ {}
+#pragma omp target has_device_addr(rh) // OK
+ {}
+#pragma omp target has_device_addr(k,i,j) // OK
+ {}
+#pragma omp target has_device_addr(d) // OK
+ {}
+#pragma omp target has_device_addr(da) // OK
+ {}
+#pragma omp target map(ps) has_device_addr(ps) // expected-error{{variable already marked as mapped in current construct}} expected-note{{used here}}
+ {}
+#pragma omp target has_device_addr(ps) map(ps) // expected-error{{variable already marked as mapped in current construct}} expected-note{{used here}}
+ {}
+#pragma omp target map(ps->a) has_device_addr(ps) // expected-error{{variable already marked as mapped in current construct}} expected-note{{used here}}
+ {}
+#pragma omp target has_device_addr(ps) map(ps->a) // expected-error{{pointer cannot be mapped along with a section derived from itself}} expected-note{{used here}}
+ {}
+#pragma omp target has_device_addr(ps) firstprivate(ps) // expected-error{{firstprivate variable cannot be in a has_device_addr clause in '#pragma omp target' directive}}
+ {}
+#pragma omp target firstprivate(ps) has_device_addr(ps) // expected-error{{firstprivate variable cannot be in a has_device_addr clause in '#pragma omp target' directive}} expected-note{{defined as firstprivate}}
+ {}
+#pragma omp target has_device_addr(ps) private(ps) // expected-error{{private variable cannot be in a has_device_addr clause in '#pragma omp target' directive}}
+ {}
+#pragma omp target private(ps) has_device_addr(ps) // expected-error{{private variable cannot be in a has_device_addr clause in '#pragma omp target' directive}} expected-note{{defined as private}}
+ {}
+ return tmain<int, 3>(argc);
+}
diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp
index 029adc6c6208a..011d9f2fd7280 100644
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@@ -2573,6 +2573,10 @@ void OMPClauseEnqueue::VisitOMPIsDevicePtrClause(
const OMPIsDevicePtrClause *C) {
VisitOMPClauseList(C);
}
+void OMPClauseEnqueue::VisitOMPHasDeviceAddrClause(
+ const OMPHasDeviceAddrClause *C) {
+ VisitOMPClauseList(C);
+}
void OMPClauseEnqueue::VisitOMPNontemporalClause(
const OMPNontemporalClause *C) {
VisitOMPClauseList(C);
diff --git a/flang/lib/Semantics/check-omp-structure.cpp b/flang/lib/Semantics/check-omp-structure.cpp
index 88d2a6d399908..2f8da1120a900 100644
--- a/flang/lib/Semantics/check-omp-structure.cpp
+++ b/flang/lib/Semantics/check-omp-structure.cpp
@@ -1683,6 +1683,7 @@ CHECK_SIMPLE_CLAUSE(Threadprivate, OMPC_threadprivate)
CHECK_SIMPLE_CLAUSE(Threads, OMPC_threads)
CHECK_SIMPLE_CLAUSE(Inbranch, OMPC_inbranch)
CHECK_SIMPLE_CLAUSE(IsDevicePtr, OMPC_is_device_ptr)
+CHECK_SIMPLE_CLAUSE(HasDeviceAddr, OMPC_has_device_addr)
CHECK_SIMPLE_CLAUSE(Link, OMPC_link)
CHECK_SIMPLE_CLAUSE(Indirect, OMPC_indirect)
CHECK_SIMPLE_CLAUSE(Mergeable, OMPC_mergeable)
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMP.td b/llvm/include/llvm/Frontend/OpenMP/OMP.td
index 053859bf361c8..e5a1dd3931247 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMP.td
+++ b/llvm/include/llvm/Frontend/OpenMP/OMP.td
@@ -253,6 +253,11 @@ def OMPC_IsDevicePtr : Clause<"is_device_ptr"> {
let flangClass = "Name";
let isValueList = true;
}
+def OMPC_HasDeviceAddr : Clause<"has_device_addr"> {
+ let clangClass = "OMPHasDeviceAddrClause";
+ let flangClass = "Name";
+ let isValueList = true;
+}
def OMPC_TaskReduction : Clause<"task_reduction"> {
let clangClass = "OMPTaskReductionClause";
let flangClass = "OmpReductionClause";
@@ -556,6 +561,7 @@ def OMP_Target : Directive<"target"> {
VersionedClause<OMPC_Depend>,
VersionedClause<OMPC_FirstPrivate>,
VersionedClause<OMPC_IsDevicePtr>,
+ VersionedClause<OMPC_HasDeviceAddr, 51>,
VersionedClause<OMPC_Reduction>,
VersionedClause<OMPC_Allocate>,
VersionedClause<OMPC_UsesAllocators, 50>
@@ -652,6 +658,7 @@ def OMP_TargetParallel : Directive<"target parallel"> {
VersionedClause<OMPC_Shared>,
VersionedClause<OMPC_Reduction>,
VersionedClause<OMPC_IsDevicePtr>,
+ VersionedClause<OMPC_HasDeviceAddr, 51>,
VersionedClause<OMPC_Allocate>,
VersionedClause<OMPC_UsesAllocators, 50>
];
@@ -684,6 +691,7 @@ def OMP_TargetParallelFor : Directive<"target parallel for"> {
VersionedClause<OMPC_Ordered>,
VersionedClause<OMPC_Linear>,
VersionedClause<OMPC_IsDevicePtr>,
+ VersionedClause<OMPC_HasDeviceAddr, 51>,
VersionedClause<OMPC_Allocate>,
VersionedClause<OMPC_Order, 50>,
VersionedClause<OMPC_UsesAllocators, 50>
@@ -700,6 +708,7 @@ def OMP_TargetParallelDo : Directive<"target parallel do"> {
VersionedClause<OMPC_Reduction>,
VersionedClause<OMPC_Linear>,
VersionedClause<OMPC_IsDevicePtr>,
+ VersionedClause<OMPC_HasDeviceAddr, 51>,
VersionedClause<OMPC_Allocator>,
VersionedClause<OMPC_Order>,
VersionedClause<OMPC_UsesAllocators>,
@@ -1133,6 +1142,7 @@ def OMP_TargetParallelForSimd : Directive<"target parallel for simd"> {
VersionedClause<OMPC_SimdLen>,
VersionedClause<OMPC_Aligned>,
VersionedClause<OMPC_IsDevicePtr>,
+ VersionedClause<OMPC_HasDeviceAddr, 51>,
VersionedClause<OMPC_Allocate>,
VersionedClause<OMPC_NonTemporal, 50>,
VersionedClause<OMPC_Order, 50>,
@@ -1163,6 +1173,7 @@ def OMP_TargetParallelDoSimd : Directive<"target parallel do simd"> {
VersionedClause<OMPC_SimdLen>,
VersionedClause<OMPC_Aligned>,
VersionedClause<OMPC_IsDevicePtr>,
+ VersionedClause<OMPC_HasDeviceAddr, 51>,
VersionedClause<OMPC_Allocate>,
VersionedClause<OMPC_NonTemporal>,
VersionedClause<OMPC_Order>,
@@ -1176,6 +1187,7 @@ def OMP_TargetSimd : Directive<"target simd"> {
VersionedClause<OMPC_Depend>,
VersionedClause<OMPC_FirstPrivate>,
VersionedClause<OMPC_IsDevicePtr>,
+ VersionedClause<OMPC_HasDeviceAddr, 51>,
VersionedClause<OMPC_LastPrivate>,
VersionedClause<OMPC_Linear>,
VersionedClause<OMPC_Map>,
@@ -1349,6 +1361,7 @@ def OMP_TargetTeams : Directive<"target teams"> {
VersionedClause<OMPC_Depend>,
VersionedClause<OMPC_FirstPrivate>,
VersionedClause<OMPC_IsDevicePtr>,
+ VersionedClause<OMPC_HasDeviceAddr, 51>,
VersionedClause<OMPC_Reduction>,
VersionedClause<OMPC_Allocate>,
VersionedClause<OMPC_UsesAllocators, 50>,
@@ -1372,6 +1385,7 @@ def OMP_TargetTeamsDistribute : Directive<"target teams distribute"> {
VersionedClause<OMPC_Depend>,
VersionedClause<OMPC_FirstPrivate>,
VersionedClause<OMPC_IsDevicePtr>,
+ VersionedClause<OMPC_HasDeviceAddr, 51>,
VersionedClause<OMPC_Reduction>,
VersionedClause<OMPC_Allocate>,
VersionedClause<OMPC_UsesAllocators, 50>,
@@ -1402,6 +1416,7 @@ def OMP_TargetTeamsDistributeParallelFor :
VersionedClause<OMPC_DefaultMap>,
VersionedClause<OMPC_FirstPrivate>,
VersionedClause<OMPC_IsDevicePtr>,
+ VersionedClause<OMPC_HasDeviceAddr, 51>,
VersionedClause<OMPC_Default>,
VersionedClause<OMPC_Shared>,
VersionedClause<OMPC_Reduction>,
@@ -1427,6 +1442,7 @@ def OMP_TargetTeamsDistributeParallelDo :
VersionedClause<OMPC_Depend>,
VersionedClause<OMPC_FirstPrivate>,
VersionedClause<OMPC_IsDevicePtr>,
+ VersionedClause<OMPC_HasDeviceAddr, 51>,
VersionedClause<OMPC_Reduction>,
VersionedClause<OMPC_Allocate>,
VersionedClause<OMPC_UsesAllocators>,
@@ -1463,6 +1479,7 @@ def OMP_TargetTeamsDistributeParallelForSimd :
VersionedClause<OMPC_DefaultMap>,
VersionedClause<OMPC_FirstPrivate>,
VersionedClause<OMPC_IsDevicePtr>,
+ VersionedClause<OMPC_HasDeviceAddr, 51>,
VersionedClause<OMPC_Default>,
VersionedClause<OMPC_Shared>,
VersionedClause<OMPC_Reduction>,
@@ -1492,6 +1509,7 @@ def OMP_TargetTeamsDistributeParallelDoSimd :
VersionedClause<OMPC_Depend>,
VersionedClause<OMPC_FirstPrivate>,
VersionedClause<OMPC_IsDevicePtr>,
+ VersionedClause<OMPC_HasDeviceAddr, 51>,
VersionedClause<OMPC_Reduction>,
VersionedClause<OMPC_Allocate>,
VersionedClause<OMPC_UsesAllocators>,
@@ -1530,6 +1548,7 @@ def OMP_TargetTeamsDistributeSimd :
VersionedClause<OMPC_FirstPrivate>,
VersionedClause<OMPC_If>,
VersionedClause<OMPC_IsDevicePtr>,
+ VersionedClause<OMPC_HasDeviceAddr, 51>,
VersionedClause<OMPC_LastPrivate>,
VersionedClause<OMPC_Linear>,
VersionedClause<OMPC_Map>,
@@ -1741,6 +1760,7 @@ def OMP_dispatch : Directive<"dispatch"> {
let allowedClauses = [
VersionedClause<OMPC_Device>,
VersionedClause<OMPC_IsDevicePtr>,
+ VersionedClause<OMPC_HasDeviceAddr, 51>,
VersionedClause<OMPC_NoWait>,
VersionedClause<OMPC_Depend>,
VersionedClause<OMPC_Novariants>,
@@ -1790,6 +1810,7 @@ def OMP_target_teams_loop : Directive<"target teams loop"> {
VersionedClause<OMPC_Device>,
VersionedClause<OMPC_FirstPrivate>,
VersionedClause<OMPC_IsDevicePtr>,
+ VersionedClause<OMPC_HasDeviceAddr, 51>,
VersionedClause<OMPC_LastPrivate>,
VersionedClause<OMPC_Map>,
VersionedClause<OMPC_Private>,
@@ -1836,6 +1857,7 @@ def OMP_target_parallel_loop : Directive<"target parallel loop"> {
VersionedClause<OMPC_Device>,
VersionedClause<OMPC_FirstPrivate>,
VersionedClause<OMPC_IsDevicePtr>,
+ VersionedClause<OMPC_HasDeviceAddr, 51>,
VersionedClause<OMPC_LastPrivate>,
VersionedClause<OMPC_Map>,
VersionedClause<OMPC_Private>,
More information about the cfe-commits
mailing list